Bart/bullet3
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

Accelerate GPU raycaster with PLBVH.

Browse Source
This commit is contained in:
Jackson Lee
2014-02-23 20:40:58 -08:00
parent e955192971
commit e4fbd5332d
10 changed files with 732 additions and 33 deletions
Download Patch File Download Diff File Expand all files Collapse all files

78
src/Bullet3OpenCL/BroadphaseCollision/b3GpuParallelLinearBvh.h
View File

@@ -18,6 +18,7 @@ subject to the following restrictions:
#include "Bullet3OpenCL/BroadphaseCollision/b3SapAabb.h"
#include "Bullet3Common/shared/b3Int2.h"
#include "Bullet3Common/shared/b3Int4.h"
#include "Bullet3Collision/NarrowPhaseCollision/b3RaycastInfo.h"
#include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h"
#include "Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h"
@@ -59,6 +60,7 @@ class b3GpuParallelLinearBvh
cl_kernel m_determineInternalNodeAabbsKernel;
cl_kernel m_plbvhCalculateOverlappingPairsKernel;
cl_kernel m_plbvhRayTraverseKernel;
b3FillCL m_fill;
b3RadixSort32CL m_radixSorter;
@@ -79,6 +81,7 @@ class b3GpuParallelLinearBvh
b3OpenCLArray<int> m_leafNodeParentNodes;
b3OpenCLArray<b3SortData> m_mortonCodesAndAabbIndicies; //m_key = morton code, m_value == aabb index
b3OpenCLArray<b3SapAabb> m_mergedAabb;
b3OpenCLArray<b3SapAabb> m_leafNodeAabbs;
public:
b3GpuParallelLinearBvh(cl_context context, cl_device_id device, cl_command_queue queue) :
@@ -94,7 +97,8 @@ public:
m_internalNodeParentNodes(context, queue),
m_leafNodeParentNodes(context, queue),
m_mortonCodesAndAabbIndicies(context, queue),
m_mergedAabb(context, queue)
m_mergedAabb(context, queue),
m_leafNodeAabbs(context, queue)
{
const char CL_PROGRAM_PATH[] = "src/Bullet3OpenCL/BroadphaseCollision/kernels/parallelLinearBvh.cl";
@@ -115,6 +119,8 @@ public:
m_plbvhCalculateOverlappingPairsKernel = b3OpenCLUtils::compileCLKernelFromString( context, device, kernelSource, "plbvhCalculateOverlappingPairs", &error, m_parallelLinearBvhProgram, additionalMacros );
b3Assert(m_plbvhCalculateOverlappingPairsKernel);
m_plbvhRayTraverseKernel = b3OpenCLUtils::compileCLKernelFromString( context, device, kernelSource, "plbvhRayTraverse", &error, m_parallelLinearBvhProgram, additionalMacros );
b3Assert(m_plbvhRayTraverseKernel);
}
virtual ~b3GpuParallelLinearBvh()
@@ -125,6 +131,7 @@ public:
clReleaseKernel(m_determineInternalNodeAabbsKernel);
clReleaseKernel(m_plbvhCalculateOverlappingPairsKernel);
clReleaseKernel(m_plbvhRayTraverseKernel);
clReleaseProgram(m_parallelLinearBvhProgram);
}
@@ -148,8 +155,12 @@ public:
m_leafNodeParentNodes.resize(numLeaves);
m_mortonCodesAndAabbIndicies.resize(numLeaves);
m_mergedAabb.resize(numLeaves);
m_leafNodeAabbs.resize(numLeaves);
}
//
m_leafNodeAabbs.copyFromOpenCLArray(worldSpaceAabbs);
//Determine number of levels in the binary tree( numLevels = ceil( log2(numLeaves) ) )
//The number of levels is equivalent to the number of bits needed to uniquely identify each node(including both internal and leaf nodes)
int numLevels = 0;
@@ -168,7 +179,7 @@ public:
if(0) printf("numLeaves, numLevels, mostSignificantBit: %d, %d, %d \n", numLeaves, numLevels, mostSignificantBit);
}
//Determine number of nodes per level, use prefix sum to get offsets of each level, and send to GPU
//Determine number of internal nodes per level, use prefix sum to get offsets of each level, and send to GPU
{
B3_PROFILE("Determine number of nodes per level");
@@ -329,7 +340,7 @@ public:
}
//For each internal node, check children to get its AABB; start from the
//last level and move towards the root
//last level, which contains the leaves, and move towards the root
{
B3_PROFILE("Set AABBs");
@@ -416,10 +427,12 @@ public:
}
}
//Max number of pairs is out_overlappingPairs.size()
//If the number of overlapping pairs is < out_overlappingPairs.size(), the array is resized
void calculateOverlappingPairs(const b3OpenCLArray<b3SapAabb>& worldSpaceAabbs,
b3OpenCLArray<int>& out_numPairs, b3OpenCLArray<b3Int4>& out_overlappingPairs)
///b3GpuParallelLinearBvh::build() must be called before this function. calculateOverlappingPairs() uses
///the worldSpaceAabbs parameter of b3GpuParallelLinearBvh::build() as the query AABBs.
///@param out_numPairs If number of pairs exceeds the max number of pairs, this is clamped to the max number.
///@param out_overlappingPairs The size() of this array is used to determine the max number of pairs.
///If the number of overlapping pairs is < out_overlappingPairs.size(), out_overlappingPairs is resized.
void calculateOverlappingPairs(b3OpenCLArray<int>& out_numPairs, b3OpenCLArray<b3Int4>& out_overlappingPairs)
{
b3Assert( out_numPairs.size() == 1 );
@@ -431,11 +444,11 @@ public:
{
B3_PROFILE("PLBVH calculateOverlappingPairs");
int numQueryAabbs = worldSpaceAabbs.size();
int numQueryAabbs = m_leafNodeAabbs.size();
b3BufferInfoCL bufferInfo[] =
{
b3BufferInfoCL( worldSpaceAabbs.getBufferCL() ),
b3BufferInfoCL( m_leafNodeAabbs.getBufferCL() ),
b3BufferInfoCL( m_internalNodeChildNodes.getBufferCL() ),
b3BufferInfoCL( m_internalNodeAabbs.getBufferCL() ),
@@ -468,6 +481,53 @@ public:
out_overlappingPairs.resize(numPairs);
}
///@param out_numRigidRayPairs Array of length 1; contains the number of detected ray-rigid AABB intersections;
///this value may be greater than out_rayRigidPairs.size() if out_rayRigidPairs is not large enough.
///@param out_rayRigidPairs Contains an array of rays intersecting rigid AABBs; x == ray index, y == rigid body index.
///If the size of this array is insufficient to hold all ray-rigid AABB intersections, additional intersections are discarded.
void testRaysAgainstBvhAabbs(const b3OpenCLArray<b3RayInfo>& rays,
b3OpenCLArray<int>& out_numRayRigidPairs, b3OpenCLArray<b3Int2>& out_rayRigidPairs)
{
B3_PROFILE("PLBVH testRaysAgainstBvhAabbs()");
int numRays = rays.size();
int maxRayRigidPairs = out_rayRigidPairs.size();
int reset = 0;
out_numRayRigidPairs.copyFromHostPointer(&reset, 1);
b3BufferInfoCL bufferInfo[] =
{
b3BufferInfoCL( m_leafNodeAabbs.getBufferCL() ),
b3BufferInfoCL( m_internalNodeChildNodes.getBufferCL() ),
b3BufferInfoCL( m_internalNodeAabbs.getBufferCL() ),
b3BufferInfoCL( m_internalNodeLeafIndexRanges.getBufferCL() ),
b3BufferInfoCL( m_mortonCodesAndAabbIndicies.getBufferCL() ),
b3BufferInfoCL( rays.getBufferCL() ),
b3BufferInfoCL( out_numRayRigidPairs.getBufferCL() ),
b3BufferInfoCL( out_rayRigidPairs.getBufferCL() )
};
b3LauncherCL launcher(m_queue, m_plbvhRayTraverseKernel, "m_plbvhRayTraverseKernel");
launcher.setBuffers( bufferInfo, sizeof(bufferInfo)/sizeof(b3BufferInfoCL) );
launcher.setConst(maxRayRigidPairs);
launcher.setConst(numRays);
launcher.launch1D(numRays);
clFinish(m_queue);
//
int numRayRigidPairs = -1;
out_numRayRigidPairs.copyToHostPointer(&numRayRigidPairs, 1);
if(numRayRigidPairs > maxRayRigidPairs)
b3Error("Error running out of rayRigid pairs: numRayRigidPairs = %d, maxRayRigidPairs = %d.\n", numRayRigidPairs, maxRayRigidPairs);
}
};
#endif

2
src/Bullet3OpenCL/BroadphaseCollision/b3GpuParallelLinearBvhBroadphase.h
View File

@@ -61,7 +61,7 @@ public:
//
m_overlappingPairsGpu.resize(maxPairs);
m_plbvh.calculateOverlappingPairs(m_aabbsGpu, m_tempNumPairs, m_overlappingPairsGpu);
m_plbvh.calculateOverlappingPairs(m_tempNumPairs, m_overlappingPairsGpu);
}
virtual void calculateOverlappingPairsHost(int maxPairs)
{

159
src/Bullet3OpenCL/BroadphaseCollision/kernels/parallelLinearBvh.cl
View File

@@ -15,6 +15,7 @@ typedef float b3Scalar;
typedef float4 b3Vector3;
#define b3Max max
#define b3Min min
#define b3Sqrt sqrt
typedef struct
{
@@ -388,3 +389,161 @@ __kernel void plbvhCalculateOverlappingPairs(__global b3AabbCL* rigidAabbs,
}
}
//From rayCastKernels.cl
typedef struct
{
float4 m_from;
float4 m_to;
} b3RayInfo;
typedef struct
{
float m_hitFraction;
int m_hitResult0;
int m_hitResult1;
int m_hitResult2;
float4 m_hitPoint;
float4 m_hitNormal;
} b3RayHit;
//From rayCastKernels.cl
b3Vector3 b3Vector3_normalize(b3Vector3 v)
{
b3Vector3 normal = (b3Vector3){v.x, v.y, v.z, 0.f};
return normalize(normal); //OpenCL normalize == vector4 normalize
}
b3Scalar b3Vector3_length2(b3Vector3 v) { return v.x*v.x + v.y*v.y + v.z*v.z; }
b3Scalar b3Vector3_dot(b3Vector3 a, b3Vector3 b) { return a.x*b.x + a.y*b.y + a.z*b.z; }
/**
int rayIntersectsAabb_optimized(b3Vector3 rayFrom, b3Vector3 rayTo, b3Vector3 rayNormalizedDirection, b3AabbCL aabb)
{
// not functional -- need to fix
//aabb is considered as 3 pairs of 2 planes( {x_min, x_max}, {y_min, y_max}, {z_min, z_max} )
//t_min is the first intersection, t_max is the second intersection
b3Vector3 inverseRayDirection = (b3Vector3){1.0f, 1.0f, 1.0f, 0.0f} / rayNormalizedDirection;
int4 sign = isless( inverseRayDirection, (b3Vector3){0.0f, 0.0f, 0.0f, 0.0f} ); //isless(x,y) returns (x < y)
//select(b, a, condition) == condition ? a : b
b3Vector3 t_min = ( select(aabb.m_min, aabb.m_max, sign) - rayFrom ) * inverseRayDirection;
b3Vector3 t_max = ( select(aabb.m_min, aabb.m_max, (int4){1,1,1,1} - sign) - rayFrom ) * inverseRayDirection;
b3Scalar t_min_final = 0.0f;
b3Scalar t_max_final = b3Sqrt( b3Vector3_length2(rayTo - rayFrom) );
//Must use fmin()/fmax(); if one of the parameters is NaN, then the parameter that is not NaN is returned.
//Behavior of min()/max() with NaNs is undefined. (See OpenCL Specification 1.2 [6.12.2] and [6.12.4])
//Since the innermost fmin()/fmax() is always not NaN, this should never return NaN
t_min_final = fmax( t_min.z, fmax(t_min.y, fmax(t_min.x, t_min_final)) );
t_max_final = fmin( t_max.z, fmin(t_max.y, fmin(t_max.x, t_max_final)) );
return (t_min_final <= t_max_final);
}
**/
void rayPlanePairTest(b3Scalar rayStart, b3Scalar rayNormalizedDirection,
b3Scalar planeMin, b3Scalar planeMax,
b3Scalar* out_t_min, b3Scalar* out_t_max)
{
if(rayNormalizedDirection < 0.0f)
{
//max is closer, min is farther
*out_t_min = (planeMax - rayStart) / rayNormalizedDirection;
*out_t_max = (planeMin - rayStart) / rayNormalizedDirection;
}
else
{
//min is closer, max is farther
*out_t_min = (planeMin - rayStart) / rayNormalizedDirection;
*out_t_max = (planeMax - rayStart) / rayNormalizedDirection;
}
}
int rayIntersectsAabb(b3Vector3 rayFrom, b3Vector3 rayTo, b3Vector3 rayNormalizedDirection, b3AabbCL aabb)
{
b3Scalar t_min_x, t_min_y, t_min_z;
b3Scalar t_max_x, t_max_y, t_max_z;
rayPlanePairTest(rayFrom.x, rayNormalizedDirection.x, aabb.m_min.x, aabb.m_max.x, &t_min_x, &t_max_x);
rayPlanePairTest(rayFrom.y, rayNormalizedDirection.y, aabb.m_min.y, aabb.m_max.y, &t_min_y, &t_max_y);
rayPlanePairTest(rayFrom.z, rayNormalizedDirection.z, aabb.m_min.z, aabb.m_max.z, &t_min_z, &t_max_z);
b3Scalar t_min_final = 0.0f;
b3Scalar t_max_final = b3Sqrt( b3Vector3_length2(rayTo - rayFrom) );
t_min_final = fmax( t_min_z, fmax(t_min_y, fmax(t_min_x, t_min_final)) );
t_max_final = fmin( t_max_z, fmin(t_max_y, fmin(t_max_x, t_max_final)) );
return (t_min_final <= t_max_final);
}
__kernel void plbvhRayTraverse(__global b3AabbCL* rigidAabbs,
__global int2* internalNodeChildIndices,
__global b3AabbCL* internalNodeAabbs,
__global int2* internalNodeLeafIndexRanges,
__global SortDataCL* mortonCodesAndAabbIndices,
__global b3RayInfo* rays,
__global int* out_numRayRigidPairs,
__global int2* out_rayRigidPairs,
int maxRayRigidPairs, int numRays)
{
int rayIndex = get_global_id(0);
if(rayIndex >= numRays) return;
b3Vector3 rayFrom = rays[rayIndex].m_from;
b3Vector3 rayTo = rays[rayIndex].m_to;
b3Vector3 rayNormalizedDirection = b3Vector3_normalize(rays[rayIndex].m_to - rays[rayIndex].m_from);
int stack[B3_PLVBH_TRAVERSE_MAX_STACK_SIZE];
//Starting by placing only the root node index, 0, in the stack causes it to be detected as a leaf node(see isLeafNode() in loop)
int stackSize = 2;
stack[0] = internalNodeChildIndices[B3_PLBVH_ROOT_NODE_INDEX].x;
stack[1] = internalNodeChildIndices[B3_PLBVH_ROOT_NODE_INDEX].y;
while(stackSize)
{
int internalOrLeafNodeIndex = stack[ stackSize - 1 ];
--stackSize;
int isLeaf = isLeafNode(internalOrLeafNodeIndex); //Internal node if false
int bvhNodeIndex = getIndexWithInternalNodeMarkerRemoved(internalOrLeafNodeIndex);
//bvhRigidIndex is not used if internal node
int bvhRigidIndex = (isLeaf) ? mortonCodesAndAabbIndices[bvhNodeIndex].m_value : -1;
b3AabbCL bvhNodeAabb = (isLeaf) ? rigidAabbs[bvhRigidIndex] : internalNodeAabbs[bvhNodeIndex];
if( rayIntersectsAabb(rayFrom, rayTo, rayNormalizedDirection, bvhNodeAabb) )
{
if(isLeaf)
{
int2 rayRigidPair;
rayRigidPair.x = rayIndex;
rayRigidPair.y = rigidAabbs[bvhRigidIndex].m_minIndices[3];
int pairIndex = atomic_inc(out_numRayRigidPairs);
if(pairIndex < maxRayRigidPairs) out_rayRigidPairs[pairIndex] = rayRigidPair;
}
if(!isLeaf) //Internal node
{
if(stackSize + 2 > B3_PLVBH_TRAVERSE_MAX_STACK_SIZE)
{
//Error
}
else
{
stack[ stackSize++ ] = internalNodeChildIndices[bvhNodeIndex].x;
stack[ stackSize++ ] = internalNodeChildIndices[bvhNodeIndex].y;
}
}
}
}
}

148
src/Bullet3OpenCL/BroadphaseCollision/kernels/parallelLinearBvhKernels.h
View File

@@ -15,6 +15,7 @@ static const char* parallelLinearBvhCL= \
"typedef float4 b3Vector3;\n"
"#define b3Max max\n"
"#define b3Min min\n"
"#define b3Sqrt sqrt\n"
"typedef struct\n"
"{\n"
" unsigned int m_key;\n"
@@ -372,4 +373,151 @@ static const char* parallelLinearBvhCL= \
" \n"
" }\n"
"}\n"
"//From rayCastKernels.cl\n"
"typedef struct\n"
"{\n"
" float4 m_from;\n"
" float4 m_to;\n"
"} b3RayInfo;\n"
"typedef struct\n"
"{\n"
" float m_hitFraction;\n"
" int m_hitResult0;\n"
" int m_hitResult1;\n"
" int m_hitResult2;\n"
" float4 m_hitPoint;\n"
" float4 m_hitNormal;\n"
"} b3RayHit;\n"
"//From rayCastKernels.cl\n"
"b3Vector3 b3Vector3_normalize(b3Vector3 v)\n"
"{\n"
" b3Vector3 normal = (b3Vector3){v.x, v.y, v.z, 0.f};\n"
" return normalize(normal); //OpenCL normalize == vector4 normalize\n"
"}\n"
"b3Scalar b3Vector3_length2(b3Vector3 v) { return v.x*v.x + v.y*v.y + v.z*v.z; }\n"
"b3Scalar b3Vector3_dot(b3Vector3 a, b3Vector3 b) { return a.x*b.x + a.y*b.y + a.z*b.z; }\n"
"/**\n"
"int rayIntersectsAabb_optimized(b3Vector3 rayFrom, b3Vector3 rayTo, b3Vector3 rayNormalizedDirection, b3AabbCL aabb)\n"
"{\n"
" // not functional -- need to fix\n"
" //aabb is considered as 3 pairs of 2 planes( {x_min, x_max}, {y_min, y_max}, {z_min, z_max} )\n"
" //t_min is the first intersection, t_max is the second intersection\n"
" b3Vector3 inverseRayDirection = (b3Vector3){1.0f, 1.0f, 1.0f, 0.0f} / rayNormalizedDirection;\n"
" int4 sign = isless( inverseRayDirection, (b3Vector3){0.0f, 0.0f, 0.0f, 0.0f} ); //isless(x,y) returns (x < y)\n"
" \n"
" //select(b, a, condition) == condition ? a : b\n"
" b3Vector3 t_min = ( select(aabb.m_min, aabb.m_max, sign) - rayFrom ) * inverseRayDirection;\n"
" b3Vector3 t_max = ( select(aabb.m_min, aabb.m_max, (int4){1,1,1,1} - sign) - rayFrom ) * inverseRayDirection;\n"
" b3Scalar t_min_final = 0.0f;\n"
" b3Scalar t_max_final = b3Sqrt( b3Vector3_length2(rayTo - rayFrom) );\n"
" \n"
" //Must use fmin()/fmax(); if one of the parameters is NaN, then the parameter that is not NaN is returned. \n"
" //Behavior of min()/max() with NaNs is undefined. (See OpenCL Specification 1.2 [6.12.2] and [6.12.4])\n"
" //Since the innermost fmin()/fmax() is always not NaN, this should never return NaN\n"
" t_min_final = fmax( t_min.z, fmax(t_min.y, fmax(t_min.x, t_min_final)) );\n"
" t_max_final = fmin( t_max.z, fmin(t_max.y, fmin(t_max.x, t_max_final)) );\n"
" \n"
" return (t_min_final <= t_max_final);\n"
"}\n"
"**/\n"
"void rayPlanePairTest(b3Scalar rayStart, b3Scalar rayNormalizedDirection,\n"
" b3Scalar planeMin, b3Scalar planeMax, \n"
" b3Scalar* out_t_min, b3Scalar* out_t_max)\n"
"{\n"
" if(rayNormalizedDirection < 0.0f)\n"
" {\n"
" //max is closer, min is farther\n"
" *out_t_min = (planeMax - rayStart) / rayNormalizedDirection;\n"
" *out_t_max = (planeMin - rayStart) / rayNormalizedDirection;\n"
" }\n"
" else\n"
" {\n"
" //min is closer, max is farther\n"
" *out_t_min = (planeMin - rayStart) / rayNormalizedDirection;\n"
" *out_t_max = (planeMax - rayStart) / rayNormalizedDirection;\n"
" }\n"
"}\n"
"int rayIntersectsAabb(b3Vector3 rayFrom, b3Vector3 rayTo, b3Vector3 rayNormalizedDirection, b3AabbCL aabb)\n"
"{\n"
" b3Scalar t_min_x, t_min_y, t_min_z;\n"
" b3Scalar t_max_x, t_max_y, t_max_z;\n"
" \n"
" rayPlanePairTest(rayFrom.x, rayNormalizedDirection.x, aabb.m_min.x, aabb.m_max.x, &t_min_x, &t_max_x);\n"
" rayPlanePairTest(rayFrom.y, rayNormalizedDirection.y, aabb.m_min.y, aabb.m_max.y, &t_min_y, &t_max_y);\n"
" rayPlanePairTest(rayFrom.z, rayNormalizedDirection.z, aabb.m_min.z, aabb.m_max.z, &t_min_z, &t_max_z);\n"
" \n"
" b3Scalar t_min_final = 0.0f;\n"
" b3Scalar t_max_final = b3Sqrt( b3Vector3_length2(rayTo - rayFrom) );\n"
" \n"
" t_min_final = fmax( t_min_z, fmax(t_min_y, fmax(t_min_x, t_min_final)) );\n"
" t_max_final = fmin( t_max_z, fmin(t_max_y, fmin(t_max_x, t_max_final)) );\n"
" \n"
" return (t_min_final <= t_max_final);\n"
"}\n"
"__kernel void plbvhRayTraverse(__global b3AabbCL* rigidAabbs,\n"
" __global int2* internalNodeChildIndices, \n"
" __global b3AabbCL* internalNodeAabbs,\n"
" __global int2* internalNodeLeafIndexRanges,\n"
" __global SortDataCL* mortonCodesAndAabbIndices,\n"
" \n"
" __global b3RayInfo* rays,\n"
" \n"
" __global int* out_numRayRigidPairs, \n"
" __global int2* out_rayRigidPairs,\n"
" int maxRayRigidPairs, int numRays)\n"
"{\n"
" int rayIndex = get_global_id(0);\n"
" if(rayIndex >= numRays) return;\n"
" \n"
" b3Vector3 rayFrom = rays[rayIndex].m_from;\n"
" b3Vector3 rayTo = rays[rayIndex].m_to;\n"
" b3Vector3 rayNormalizedDirection = b3Vector3_normalize(rays[rayIndex].m_to - rays[rayIndex].m_from);\n"
" \n"
" int stack[B3_PLVBH_TRAVERSE_MAX_STACK_SIZE];\n"
" \n"
" //Starting by placing only the root node index, 0, in the stack causes it to be detected as a leaf node(see isLeafNode() in loop)\n"
" int stackSize = 2;\n"
" stack[0] = internalNodeChildIndices[B3_PLBVH_ROOT_NODE_INDEX].x;\n"
" stack[1] = internalNodeChildIndices[B3_PLBVH_ROOT_NODE_INDEX].y;\n"
" \n"
" while(stackSize)\n"
" {\n"
" int internalOrLeafNodeIndex = stack[ stackSize - 1 ];\n"
" --stackSize;\n"
" \n"
" int isLeaf = isLeafNode(internalOrLeafNodeIndex); //Internal node if false\n"
" int bvhNodeIndex = getIndexWithInternalNodeMarkerRemoved(internalOrLeafNodeIndex);\n"
" \n"
" //bvhRigidIndex is not used if internal node\n"
" int bvhRigidIndex = (isLeaf) ? mortonCodesAndAabbIndices[bvhNodeIndex].m_value : -1;\n"
" \n"
" b3AabbCL bvhNodeAabb = (isLeaf) ? rigidAabbs[bvhRigidIndex] : internalNodeAabbs[bvhNodeIndex];\n"
" \n"
" if( rayIntersectsAabb(rayFrom, rayTo, rayNormalizedDirection, bvhNodeAabb) )\n"
" {\n"
" if(isLeaf)\n"
" {\n"
" int2 rayRigidPair;\n"
" rayRigidPair.x = rayIndex;\n"
" rayRigidPair.y = rigidAabbs[bvhRigidIndex].m_minIndices[3];\n"
" \n"
" int pairIndex = atomic_inc(out_numRayRigidPairs);\n"
" if(pairIndex < maxRayRigidPairs) out_rayRigidPairs[pairIndex] = rayRigidPair;\n"
" }\n"
" \n"
" if(!isLeaf) //Internal node\n"
" {\n"
" if(stackSize + 2 > B3_PLVBH_TRAVERSE_MAX_STACK_SIZE)\n"
" {\n"
" //Error\n"
" }\n"
" else\n"
" {\n"
" stack[ stackSize++ ] = internalNodeChildIndices[bvhNodeIndex].x;\n"
" stack[ stackSize++ ] = internalNodeChildIndices[bvhNodeIndex].y;\n"
" }\n"
" }\n"
" }\n"
" }\n"
"}\n"
;