OBB Tree Parallel creation

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello everybody!

I am trying to change the code of the gProximity library for my Master’s. My problems are on OBB tree creation. The code can create tight OBB for leafs nodes without problems but when it tries to create a parent based on children OBBs parts of them are not being bounded entirely. The relevant code for intern nodes creation is shown below as the class definition. Any help would be very appreciated.

__device__ __inline__ void init(OBB &bv1, OBB &bv2)

        {

                // somewhat based on code from Dave Eberly / geometrictools.com

// guess at center by taking middle

                center = f3v_add(bv1.center, bv2.center);

                center = f3v_mul1(center, 0.5f);

// ASSUME: axis = axis of BV1

                axis1 = bv1.axis1;

                axis2 = bv1.axis2;

                axis3 = bv1.axis3;

// extents as calculated from all 16 points and 3 axes

                float3 minExtent = make_float3(CUDART_NORM_HUGE_F, CUDART_NORM_HUGE_F, CUDART_NORM_HUGE_F),

                       maxExtent = make_float3(-CUDART_NORM_HUGE_F, -CUDART_NORM_HUGE_F, -CUDART_NORM_HUGE_F);

calcExtentsForOBB(bv1, minExtent, maxExtent);

                calcExtentsForOBB(bv2, minExtent, maxExtent);

//

                float3 distCenters = f3v_sub(bv2.center , bv1.center);

                float3 distCentersHalf = f3v_mul1(distCenters , 0.5);

                minExtent = f3v_sub(minExtent , distCentersHalf);

                maxExtent = f3v_add(maxExtent , distCentersHalf);

                //

// [minExtent,maxExtent] is the axis-aligned box in the coordinate system of the

                // merged box axes.  Update the current box center to be the center of

                // the new box.  Compute the extents based on the new center.

float factorX = 0.5f * (maxExtent.x + minExtent.x);

                float3 correctorX = f3v_mul1(axis1, factorX);

                center = f3v_add(correctorX, center);

                extents.x = 0.5f * (maxExtent.x - minExtent.x);

float factorY = 0.5f * (maxExtent.y + minExtent.y);

                float3 correctorY = f3v_mul1(axis2, factorY);

                center = f3v_add(correctorY, center);

                extents.y = 0.5f * (maxExtent.y - minExtent.y);

float factorZ = 0.5f * (maxExtent.z + minExtent.z);

                float3 correctorZ = f3v_mul1(axis3, factorZ);

                center = f3v_add(correctorZ, center);

                extents.z = 0.5f * (maxExtent.z - minExtent.z);

        }

__host__ __device__ __inline__ void calcExtentsForOBB(OBB &bv, float3 &minExtent, float3 &maxExtent)

        {

                // generate all 8 points of OBB

                bv.axis1 = f3v_mul1(bv.axis1, bv.extents.x);

                bv.axis2 = f3v_mul1(bv.axis2, bv.extents.y);

                bv.axis3 = f3v_mul1(bv.axis3, bv.extents.z);

float3 point, temp;

                float dist;

// 1

                temp = f3v_add(bv.axis1, bv.axis2);

                point = f3v_add(temp, bv.axis3);

                dist = f3v_dot(point, axis1);

                minExtent.x = min(minExtent.x, dist);

                maxExtent.x = max(maxExtent.x, dist);

                dist = f3v_dot(point, axis2);

                minExtent.y = min(minExtent.y, dist);

                maxExtent.y = max(maxExtent.y, dist);

                dist = f3v_dot(point, axis3);

                minExtent.z = min(minExtent.z, dist);

                maxExtent.z = max(maxExtent.z, dist);

// 2

                point = f3v_sub(temp, bv.axis3);

                dist = f3v_dot(point, axis1);

                minExtent.x = min(minExtent.x, dist);

                maxExtent.x = max(maxExtent.x, dist);

                dist = f3v_dot(point, axis2);

                minExtent.y = min(minExtent.y, dist);

                maxExtent.y = max(maxExtent.y, dist);

                dist = f3v_dot(point, axis3);

                minExtent.z = min(minExtent.z, dist);

                maxExtent.z = max(maxExtent.z, dist);

// 3

                temp = f3v_sub(bv.axis1, bv.axis2);

                point = f3v_add(temp, bv.axis3);

                dist = f3v_dot(point, axis1);

                minExtent.x = min(minExtent.x, dist);

                maxExtent.x = max(maxExtent.x, dist);

                dist = f3v_dot(point, axis2);

                minExtent.y = min(minExtent.y, dist);

                maxExtent.y = max(maxExtent.y, dist);

                dist = f3v_dot(point, axis3);

                minExtent.z = min(minExtent.z, dist);

                maxExtent.z = max(maxExtent.z, dist);

// 4

                point = f3v_sub(temp, bv.axis3);

                dist = f3v_dot(point, axis1);

                minExtent.x = min(minExtent.x, dist);

                maxExtent.x = max(maxExtent.x, dist);

                dist = f3v_dot(point, axis2);

                minExtent.y = min(minExtent.y, dist);

                maxExtent.y = max(maxExtent.y, dist);

                dist = f3v_dot(point, axis3);

                minExtent.z = min(minExtent.z, dist);

                maxExtent.z = max(maxExtent.z, dist);

// 5

                bv.axis1.x = -bv.axis1.x;

                bv.axis1.y = -bv.axis1.y;

                bv.axis1.z = -bv.axis1.z;

temp = f3v_add(bv.axis1, bv.axis2);

                point = f3v_add(temp, bv.axis3);

                dist = f3v_dot(point, axis1);

                minExtent.x = min(minExtent.x, dist);

                maxExtent.x = max(maxExtent.x, dist);

                dist = f3v_dot(point, axis2);

                minExtent.y = min(minExtent.y, dist);

                maxExtent.y = max(maxExtent.y, dist);

                dist = f3v_dot(point, axis3);

                minExtent.z = min(minExtent.z, dist);

                maxExtent.z = max(maxExtent.z, dist);

// 6

                point = f3v_sub(temp, bv.axis3);

                dist = f3v_dot(point, axis1);

                minExtent.x = min(minExtent.x, dist);

                maxExtent.x = max(maxExtent.x, dist);

                dist = f3v_dot(point, axis2);

                minExtent.y = min(minExtent.y, dist);

                maxExtent.y = max(maxExtent.y, dist);

                dist = f3v_dot(point, axis3);

                minExtent.z = min(minExtent.z, dist);

                maxExtent.z = max(maxExtent.z, dist);

// 7

                temp = f3v_sub(bv.axis1, bv.axis2);

                point = f3v_add(temp, bv.axis3);

                dist = f3v_dot(point, axis1);

                minExtent.x = min(minExtent.x, dist);

                maxExtent.x = max(maxExtent.x, dist);

                dist = f3v_dot(point, axis2);

                minExtent.y = min(minExtent.y, dist);

                maxExtent.y = max(maxExtent.y, dist);

                dist = f3v_dot(point, axis3);

                minExtent.z = min(minExtent.z, dist);

                maxExtent.z = max(maxExtent.z, dist);

// 8

                point = f3v_sub(temp, bv.axis3);

                dist = f3v_dot(point, axis1);

                minExtent.x = min(minExtent.x, dist);

                maxExtent.x = max(maxExtent.x, dist);

                dist = f3v_dot(point, axis2);

                minExtent.y = min(minExtent.y, dist);

                maxExtent.y = max(maxExtent.y, dist);

                dist = f3v_dot(point, axis3);

                minExtent.z = min(minExtent.z, dist);

                maxExtent.z = max(maxExtent.z, dist);

        }

class OBB

{

public:

        __host__ __device__ __inline__ void calcExtentsForOBB(OBB &bv, float3 &minExtent, float3 &maxExtent);

__host__ __device__ __inline__ void calcOBBPoints(float3 *result_points) const;

__device__ __inline__ void init(OBB &bv1, OBB &bv2);

// init OBB from triangle

        __device__ __inline__ void init(const float3 &v1, const float3 &v2, const float3 &v3);

__device__ __inline__ void init(const float3 &v1, const float3 &v2, const float3 &v3, const float3 &u1, const float3 &u2, const float3 &u3);

__device__ __inline__ float getSize() const;

//Draws the box (just the side faces)

        __host__ __inline__ void draw();

float3 center;          // center point of OBB

        float3 axis1,           // major axes specifying the OBB

               axis2,

               axis3;

        float3 extents;         // extents to boundary from center point,

                                                // along each axis

};

Thanks in advance,

Vinícius da Silva.

2

Just a thought, Tthis looks like it is to run on a single thread as it doesn’t appear to use threadIdx or blockIdx though they might be in code you didn’t include in the post.

3

I didn’t put the code for thread management, just the code that a thread use to set the boundaries of the OBB, where the problem was. But I’ve already figured out the solution. The code was ignoring the distance between the two child OBB centers. Because of that, it was calculating the extents based only on one center, like the two OBBs shared the same center, which was wrong. I changed the calcExtentsForOBB function to project the centers distance on the parent OBB axis and add it to the calculated extents. Now the OBBs are conservative. The new code of the function is following.

__host__ __device__ __inline__ void calcExtentsForOBB(OBB &bv, float3 &minExtent, float3 &maxExtent)

	{

		float dist;

		float3 point, temp;

		float3 centersDistance = f3v_sub(bv.center , center);

		float3 centersProjectedDistance;

		// generate all 8 points of OBB

		centersProjectedDistance.x = f3v_dot(centersDistance , axis1);

		centersProjectedDistance.y = f3v_dot(centersDistance , axis2);

		centersProjectedDistance.z = f3v_dot(centersDistance , axis3);

		bv.axis1 = f3v_mul1(bv.axis1, bv.extents.x);

		bv.axis2 = f3v_mul1(bv.axis2, bv.extents.y);

		bv.axis3 = f3v_mul1(bv.axis3, bv.extents.z);

		

		// 1

		temp = f3v_add(bv.axis1, bv.axis2);

		point = f3v_add(temp, bv.axis3);

		dist = f3v_dot(point, axis1) + centersProjectedDistance.x;

		minExtent.x = min(minExtent.x, dist);

		maxExtent.x = max(maxExtent.x, dist);

		dist = f3v_dot(point, axis2) + centersProjectedDistance.y;

		minExtent.y = min(minExtent.y, dist);

		maxExtent.y = max(maxExtent.y, dist);

		dist = f3v_dot(point, axis3) + centersProjectedDistance.z;

		minExtent.z = min(minExtent.z, dist);

		maxExtent.z = max(maxExtent.z, dist);

		// 2

		point = f3v_sub(temp, bv.axis3);

		dist = f3v_dot(point, axis1) + centersProjectedDistance.x;

		minExtent.x = min(minExtent.x, dist);

		maxExtent.x = max(maxExtent.x, dist);

		dist = f3v_dot(point, axis2) + centersProjectedDistance.y;

		minExtent.y = min(minExtent.y, dist);

		maxExtent.y = max(maxExtent.y, dist);

		dist = f3v_dot(point, axis3) + centersProjectedDistance.z;

		minExtent.z = min(minExtent.z, dist);

		maxExtent.z = max(maxExtent.z, dist);

		

		// 3

		temp = f3v_sub(bv.axis1, bv.axis2);

		point = f3v_add(temp, bv.axis3);

		dist = f3v_dot(point, axis1) + centersProjectedDistance.x;

		minExtent.x = min(minExtent.x, dist);

		maxExtent.x = max(maxExtent.x, dist);

		dist = f3v_dot(point, axis2) + centersProjectedDistance.y;

		minExtent.y = min(minExtent.y, dist);

		maxExtent.y = max(maxExtent.y, dist);

		dist = f3v_dot(point, axis3) + centersProjectedDistance.z;

		minExtent.z = min(minExtent.z, dist);

		maxExtent.z = max(maxExtent.z, dist);

		

		// 4

		point = f3v_sub(temp, bv.axis3);

		dist = f3v_dot(point, axis1) + centersProjectedDistance.x;

		minExtent.x = min(minExtent.x, dist);

		maxExtent.x = max(maxExtent.x, dist);

		dist = f3v_dot(point, axis2) + centersProjectedDistance.y;

		minExtent.y = min(minExtent.y, dist);

		maxExtent.y = max(maxExtent.y, dist);

		dist = f3v_dot(point, axis3) + centersProjectedDistance.z;

		minExtent.z = min(minExtent.z, dist);

		maxExtent.z = max(maxExtent.z, dist);

		

		// 5

		bv.axis1.x = -bv.axis1.x;

		bv.axis1.y = -bv.axis1.y;

		bv.axis1.z = -bv.axis1.z;

		

		temp = f3v_add(bv.axis1, bv.axis2);

		point = f3v_add(temp, bv.axis3);

		dist = f3v_dot(point, axis1) + centersProjectedDistance.x;

		minExtent.x = min(minExtent.x, dist);

		maxExtent.x = max(maxExtent.x, dist);

		dist = f3v_dot(point, axis2) + centersProjectedDistance.y;

		minExtent.y = min(minExtent.y, dist);

		maxExtent.y = max(maxExtent.y, dist);

		dist = f3v_dot(point, axis3) + centersProjectedDistance.z;

		minExtent.z = min(minExtent.z, dist);

		maxExtent.z = max(maxExtent.z, dist);

		

		// 6

		point = f3v_sub(temp, bv.axis3);

		dist = f3v_dot(point, axis1) + centersProjectedDistance.x;

		minExtent.x = min(minExtent.x, dist);

		maxExtent.x = max(maxExtent.x, dist);

		dist = f3v_dot(point, axis2) + centersProjectedDistance.y;

		minExtent.y = min(minExtent.y, dist);

		maxExtent.y = max(maxExtent.y, dist);

		dist = f3v_dot(point, axis3) + centersProjectedDistance.z;

		minExtent.z = min(minExtent.z, dist);

		maxExtent.z = max(maxExtent.z, dist);

		

		// 7

		temp = f3v_sub(bv.axis1, bv.axis2);

		point = f3v_add(temp, bv.axis3);

		dist = f3v_dot(point, axis1) + centersProjectedDistance.x;

		minExtent.x = min(minExtent.x, dist);

		maxExtent.x = max(maxExtent.x, dist);

		dist = f3v_dot(point, axis2) + centersProjectedDistance.y;

		minExtent.y = min(minExtent.y, dist);

		maxExtent.y = max(maxExtent.y, dist);

		dist = f3v_dot(point, axis3) + centersProjectedDistance.z;

		minExtent.z = min(minExtent.z, dist);

		maxExtent.z = max(maxExtent.z, dist);

		

		// 8

		point = f3v_sub(temp, bv.axis3);

		dist = f3v_dot(point, axis1) + centersProjectedDistance.x;

		minExtent.x = min(minExtent.x, dist);

		maxExtent.x = max(maxExtent.x, dist);

		dist = f3v_dot(point, axis2) + centersProjectedDistance.y;

		minExtent.y = min(minExtent.y, dist);

		maxExtent.y = max(maxExtent.y, dist);

		dist = f3v_dot(point, axis3) + centersProjectedDistance.z;

		minExtent.z = min(minExtent.z, dist);

		maxExtent.z = max(maxExtent.z, dist);

	}