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

make bvh kernel work on AMD GPU (and reported an OpenCL compiler bug)

Browse Source 
enable source-level OpenCL debugging of a single kernel
This commit is contained in:
erwin coumans
2013-03-20 13:54:00 -07:00
parent 39884e4665
commit 9a693fb850
6 changed files with 192 additions and 171 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
demo/gpudemo/main_opengl3core.cpp
View File

@@ -359,6 +359,8 @@ const char* g_deviceName = "blaat";
int main(int argc, char* argv[])
{
printf("main start");
CommandLineArgs args(argc,argv);

12
opencl/basic_initialize/btOpenCLUtils.cpp
View File

@@ -16,6 +16,10 @@ subject to the following restrictions:
//original author: Roman Ponomarev
//cleanup by Erwin Coumans
bool gDebugForceLoadingFromSource = false;
bool gDebugSkipLoadingBinary = false;
#include <string.h>
#ifdef _WIN32
@@ -57,8 +61,6 @@ static const char* spPlatformVendor =
#endif //_WIN32
#endif
bool gDebugForceLoadingFromSource = false;
bool gDebugSkipLoadingBinary = false;
void MyFatalBreakAPPLE( const char * errstr ,
const void * private_info ,
@@ -519,7 +521,7 @@ static const char* strip2(const char* name, const char* pattern)
return oriptr;
}
cl_program btOpenCLUtils_compileCLProgramFromString(cl_context clContext, cl_device_id device, const char* kernelSourceOrg, cl_int* pErrNum, const char* additionalMacrosArg , const char* clFileNameForCaching)
cl_program btOpenCLUtils_compileCLProgramFromString(cl_context clContext, cl_device_id device, const char* kernelSourceOrg, cl_int* pErrNum, const char* additionalMacrosArg , const char* clFileNameForCaching, bool disableBinaryCaching)
{
const char* additionalMacros = additionalMacrosArg?additionalMacrosArg:"";
@@ -530,7 +532,7 @@ cl_program btOpenCLUtils_compileCLProgramFromString(cl_context clContext, cl_dev
char binaryFileName[BT_MAX_STRING_LENGTH];
char* bla=0;
if (clFileNameForCaching && !(gDebugSkipLoadingBinary||gDebugForceLoadingFromSource) )
if (clFileNameForCaching && !(disableBinaryCaching || gDebugSkipLoadingBinary||gDebugForceLoadingFromSource) )
{
char deviceName[256];
@@ -874,7 +876,7 @@ cl_kernel btOpenCLUtils_compileCLKernelFromString(cl_context clContext, cl_devic
if (!m_cpProgram)
{
m_cpProgram = btOpenCLUtils_compileCLProgramFromString(clContext,device,kernelSource,pErrNum, additionalMacros,0);
m_cpProgram = btOpenCLUtils_compileCLProgramFromString(clContext,device,kernelSource,pErrNum, additionalMacros,0, false);
}

6
opencl/basic_initialize/btOpenCLUtils.h
View File

@@ -41,7 +41,7 @@ void btOpenCLUtils_printDeviceInfo(cl_device_id device);
cl_kernel btOpenCLUtils_compileCLKernelFromString( cl_context clContext,cl_device_id device, const char* kernelSource, const char* kernelName, cl_int* pErrNum, cl_program prog,const char* additionalMacros);
//optional
cl_program btOpenCLUtils_compileCLProgramFromString( cl_context clContext,cl_device_id device, const char* kernelSource, cl_int* pErrNum,const char* additionalMacros , const char* srcFileNameForCaching);
cl_program btOpenCLUtils_compileCLProgramFromString( cl_context clContext,cl_device_id device, const char* kernelSource, cl_int* pErrNum,const char* additionalMacros , const char* srcFileNameForCaching, bool disableBinaryCaching);
//the following optional APIs provide access using specific platform information
int btOpenCLUtils_getNumPlatforms(cl_int* pErrNum);
@@ -141,9 +141,9 @@ struct btOpenCLUtils
}
//optional
static inline cl_program compileCLProgramFromString( cl_context clContext,cl_device_id device, const char* kernelSource, cl_int* pErrNum=0,const char* additionalMacros = "" , const char* srcFileNameForCaching=0)
static inline cl_program compileCLProgramFromString( cl_context clContext,cl_device_id device, const char* kernelSource, cl_int* pErrNum=0,const char* additionalMacros = "" , const char* srcFileNameForCaching=0, bool disableBinaryCaching=false)
{
return btOpenCLUtils_compileCLProgramFromString(clContext,device, kernelSource, pErrNum,additionalMacros, srcFileNameForCaching);
return btOpenCLUtils_compileCLProgramFromString(clContext,device, kernelSource, pErrNum,additionalMacros, srcFileNameForCaching, disableBinaryCaching);
}
//the following optional APIs provide access using specific platform information

9
opencl/gpu_sat/host/ConvexHullContact.cpp
View File

@@ -111,10 +111,11 @@ m_totalContactsOut(m_context, m_queue)
if (1)
{
const char* srcBvh = bvhTraversalKernelCL;
cl_program bvhTraversalProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,srcBvh,&errNum,"","opencl/gpu_sat/kernels/bvhTraversal.cl");
//cl_program bvhTraversalProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,srcBvh,&errNum,"","opencl/gpu_sat/kernels/bvhTraversal.cl");
cl_program bvhTraversalProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,0,&errNum,"","opencl/gpu_sat/kernels/bvhTraversal.cl", true);
btAssert(errNum==CL_SUCCESS);
m_bvhTraversalKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,srcBvh, "bvhTraversalKernel",&errNum,bvhTraversalProg);
m_bvhTraversalKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,srcBvh, "bvhTraversalKernel",&errNum,bvhTraversalProg,"-g");
btAssert(errNum==CL_SUCCESS);
}
@@ -329,6 +330,8 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
btOpenCLArray<btBvhSubtreeInfo> subTreesGPU(this->m_context,this->m_queue,numSubTrees);
subTreesGPU.copyFromHost(bvhData[0]->getSubtreeInfoArray());
btVector3 bvhAabbMin = bvhData[0]->m_bvhAabbMin;
btVector3 bvhAabbMax = bvhData[0]->m_bvhAabbMax;
btVector3 bvhQuantization = bvhData[0]->m_bvhQuantization;
@@ -357,7 +360,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
triangleConvexPairsOut.resize(np);
btAlignedObjectArray<btInt4> pairsOutCPU;
triangleConvexPairsOut.copyToHost(pairsOutCPU);
clFinish(m_queue);
printf("np=%d\n", np);

167
opencl/gpu_sat/kernels/bvhTraversal.cl
View File

@@ -47,7 +47,7 @@ typedef struct
}
*/
int getTriangleIndex(__global const btQuantizedBvhNode* rootNode)
int getTriangleIndex(const btQuantizedBvhNode* rootNode)
{
unsigned int x=0;
unsigned int y = (~(x&0))<<(31-MAX_NUM_PARTS_IN_BITS);
@@ -55,13 +55,13 @@ int getTriangleIndex(__global const btQuantizedBvhNode* rootNode)
return (rootNode->m_escapeIndexOrTriangleIndex&~(y));
}
bool isLeaf(__global const btQuantizedBvhNode* rootNode)
int isLeaf(const btQuantizedBvhNode* rootNode)
{
//skipindex is negative (internal node), triangleindex >=0 (leafnode)
return (rootNode->m_escapeIndexOrTriangleIndex >= 0);
return (rootNode->m_escapeIndexOrTriangleIndex >= 0)? 1 : 0;
}
int getEscapeIndex(__global const btQuantizedBvhNode* rootNode)
int getEscapeIndex(const btQuantizedBvhNode* rootNode)
{
return -rootNode->m_escapeIndexOrTriangleIndex;
}
@@ -129,13 +129,30 @@ typedef struct
} btAabbCL;
bool testQuantizedAabbAgainstQuantizedAabb(__private const unsigned short int* aabbMin1,__private const unsigned short int* aabbMax1,__global const unsigned short int* aabbMin2,__global const unsigned short int* aabbMax2)
int testQuantizedAabbAgainstQuantizedAabb(
const unsigned short int* aabbMin1,
const unsigned short int* aabbMax1,
const unsigned short int* aabbMin2,
const unsigned short int* aabbMax2)
{
bool overlap = true;
overlap = (aabbMin1[0] > aabbMax2[0] || aabbMax1[0] < aabbMin2[0]) ? false : overlap;
overlap = (aabbMin1[2] > aabbMax2[2] || aabbMax1[2] < aabbMin2[2]) ? false : overlap;
overlap = (aabbMin1[1] > aabbMax2[1] || aabbMax1[1] < aabbMin2[1]) ? false : overlap;
return overlap;
//int overlap = 1;
if (aabbMin1[0] > aabbMax2[0])
return 0;
if (aabbMax1[0] < aabbMin2[0])
return 0;
if (aabbMin1[1] > aabbMax2[1])
return 0;
if (aabbMax1[1] < aabbMin2[1])
return 0;
if (aabbMin1[2] > aabbMax2[2])
return 0;
if (aabbMax1[2] < aabbMin2[2])
return 0;
return 1;
//overlap = ((aabbMin1[0] > aabbMax2[0]) || (aabbMax1[0] < aabbMin2[0])) ? 0 : overlap;
//overlap = ((aabbMin1[2] > aabbMax2[2]) || (aabbMax1[2] < aabbMin2[2])) ? 0 : overlap;
//overlap = ((aabbMin1[1] > aabbMax2[1]) || (aabbMax1[1] < aabbMin2[1])) ? 0 : overlap;
//return overlap;
}
@@ -176,87 +193,77 @@ __kernel void bvhTraversalKernel( __global const int2* pairs,
int numPairs,
int maxNumConcavePairsCapacity)
{
int i = get_global_id(0);
int id = get_global_id(0);
if (id>=numPairs)
return;
if (i<numPairs)
int bodyIndexA = pairs[id].x;
int bodyIndexB = pairs[id].y;
int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;
int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;
//once the broadphase avoids static-static pairs, we can remove this test
if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))
{
return;
}
if (collidables[collidableIndexA].m_shapeType!=SHAPE_CONCAVE_TRIMESH)
return;
if (collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL)
return;
unsigned short int quantizedQueryAabbMin[3];
unsigned short int quantizedQueryAabbMax[3];
quantizeWithClamp(quantizedQueryAabbMin,aabbs[bodyIndexB].m_min,false,bvhAabbMin, bvhAabbMax,bvhQuantization);
quantizeWithClamp(quantizedQueryAabbMax,aabbs[bodyIndexB].m_max,true ,bvhAabbMin, bvhAabbMax,bvhQuantization);
int bodyIndexA = pairs[i].x;
int bodyIndexB = pairs[i].y;
int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;
int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;
int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;
int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;
//once the broadphase avoids static-static pairs, we can remove this test
if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))
for (int i=0;i<numSubtreeHeaders;i++)
{
btBvhSubtreeInfo subtree = subtreeHeaders[i];
int overlap = testQuantizedAabbAgainstQuantizedAabb(quantizedQueryAabbMin,quantizedQueryAabbMax,subtree.m_quantizedAabbMin,subtree.m_quantizedAabbMax);
if (overlap != 0)
{
return;
}
if ((collidables[collidableIndexA].m_shapeType==SHAPE_CONCAVE_TRIMESH))// && (collidables[collidableIndexB].m_shapeType==SHAPE_CONVEX_HULL))
{
unsigned short int quantizedQueryAabbMin[3];
unsigned short int quantizedQueryAabbMax[3];
quantizeWithClamp(quantizedQueryAabbMin,aabbs[bodyIndexB].m_min,false,bvhAabbMin, bvhAabbMax,bvhQuantization);
quantizeWithClamp(quantizedQueryAabbMax,aabbs[bodyIndexB].m_max,true ,bvhAabbMin, bvhAabbMax,bvhQuantization);
int i;
for (i=0;i<numSubtreeHeaders;i++)
int startNodeIndex = subtree.m_rootNodeIndex;
int endNodeIndex = subtree.m_rootNodeIndex+subtree.m_subtreeSize;
int curIndex = startNodeIndex;
int escapeIndex;
int isLeafNode;
int aabbOverlap;
while (curIndex < endNodeIndex)
{
const __global btBvhSubtreeInfo* subtree = &subtreeHeaders[i];
//PCK: unsigned instead of bool
unsigned overlap = testQuantizedAabbAgainstQuantizedAabb(quantizedQueryAabbMin,quantizedQueryAabbMax,subtree->m_quantizedAabbMin,subtree->m_quantizedAabbMax);
if (overlap != 0)
btQuantizedBvhNode rootNode = quantizedNodes[curIndex];
aabbOverlap = testQuantizedAabbAgainstQuantizedAabb(quantizedQueryAabbMin,quantizedQueryAabbMax,rootNode.m_quantizedAabbMin,rootNode.m_quantizedAabbMax);
isLeafNode = isLeaf(&rootNode);
if (aabbOverlap)
{
int startNodeIndex = subtree->m_rootNodeIndex;
int endNodeIndex = subtree->m_rootNodeIndex+subtree->m_subtreeSize;
int curIndex = startNodeIndex;
int subTreeSize = endNodeIndex - startNodeIndex;
__global const btQuantizedBvhNode* rootNode = &quantizedNodes[startNodeIndex];
int escapeIndex;
bool isLeafNode;
unsigned aabbOverlap;
while (curIndex < endNodeIndex)
if (isLeafNode)
{
aabbOverlap = testQuantizedAabbAgainstQuantizedAabb(quantizedQueryAabbMin,quantizedQueryAabbMax,rootNode->m_quantizedAabbMin,rootNode->m_quantizedAabbMax);
isLeafNode = isLeaf(rootNode);
if (isLeafNode && aabbOverlap)
int triangleIndex = getTriangleIndex(&rootNode);
int pairIdx = atomic_inc(numConcavePairsOut);
if (pairIdx<maxNumConcavePairsCapacity)
{
//do your thing! nodeCallback->processNode(rootNode->getPartId(),rootNode->getTriangleIndex());
int triangleIndex = getTriangleIndex(rootNode);
int pairIdx = atomic_inc(numConcavePairsOut);
if (pairIdx<maxNumConcavePairsCapacity)
{
//int4 newPair;
concavePairsOut[pairIdx].x = bodyIndexA;
concavePairsOut[pairIdx].y = bodyIndexB;
concavePairsOut[pairIdx].z = triangleIndex;
concavePairsOut[pairIdx].w = 3;
}
}
if ((aabbOverlap != 0) || isLeafNode)
{
rootNode++;
curIndex++;
} else
{
escapeIndex = getEscapeIndex(rootNode);
rootNode += escapeIndex;
curIndex += escapeIndex;
int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,3);
concavePairsOut[pairIdx] = newPair;
}
}
curIndex++;
} else
{
if (isLeafNode)
{
curIndex++;
} else
{
escapeIndex = getEscapeIndex(&rootNode);
curIndex += escapeIndex;
}
}
}
}//SHAPE_CONCAVE_TRIMESH
}//i<numpairs
}
}
}

167
opencl/gpu_sat/kernels/bvhTraversal.h
View File

@@ -49,7 +49,7 @@ static const char* bvhTraversalKernelCL= \
" }\n"
"*/\n"
"\n"
"int getTriangleIndex(__global const btQuantizedBvhNode* rootNode)\n"
"int getTriangleIndex(const btQuantizedBvhNode* rootNode)\n"
"{\n"
" unsigned int x=0;\n"
" unsigned int y = (~(x&0))<<(31-MAX_NUM_PARTS_IN_BITS);\n"
@@ -57,13 +57,13 @@ static const char* bvhTraversalKernelCL= \
" return (rootNode->m_escapeIndexOrTriangleIndex&~(y));\n"
"}\n"
"\n"
"bool isLeaf(__global const btQuantizedBvhNode* rootNode)\n"
"int isLeaf(const btQuantizedBvhNode* rootNode)\n"
"{\n"
" //skipindex is negative (internal node), triangleindex >=0 (leafnode)\n"
" return (rootNode->m_escapeIndexOrTriangleIndex >= 0);\n"
" return (rootNode->m_escapeIndexOrTriangleIndex >= 0)? 1 : 0;\n"
"}\n"
" \n"
"int getEscapeIndex(__global const btQuantizedBvhNode* rootNode)\n"
"int getEscapeIndex(const btQuantizedBvhNode* rootNode)\n"
"{\n"
" return -rootNode->m_escapeIndexOrTriangleIndex;\n"
"}\n"
@@ -131,13 +131,30 @@ static const char* bvhTraversalKernelCL= \
"} btAabbCL;\n"
"\n"
"\n"
"bool testQuantizedAabbAgainstQuantizedAabb(__private const unsigned short int* aabbMin1,__private const unsigned short int* aabbMax1,__global const unsigned short int* aabbMin2,__global const unsigned short int* aabbMax2)\n"
"int testQuantizedAabbAgainstQuantizedAabb(\n"
" const unsigned short int* aabbMin1,\n"
" const unsigned short int* aabbMax1,\n"
" const unsigned short int* aabbMin2,\n"
" const unsigned short int* aabbMax2)\n"
"{\n"
" bool overlap = true;\n"
" overlap = (aabbMin1[0] > aabbMax2[0] || aabbMax1[0] < aabbMin2[0]) ? false : overlap;\n"
" overlap = (aabbMin1[2] > aabbMax2[2] || aabbMax1[2] < aabbMin2[2]) ? false : overlap;\n"
" overlap = (aabbMin1[1] > aabbMax2[1] || aabbMax1[1] < aabbMin2[1]) ? false : overlap;\n"
" return overlap;\n"
" //int overlap = 1;\n"
" if (aabbMin1[0] > aabbMax2[0])\n"
" return 0;\n"
" if (aabbMax1[0] < aabbMin2[0])\n"
" return 0;\n"
" if (aabbMin1[1] > aabbMax2[1])\n"
" return 0;\n"
" if (aabbMax1[1] < aabbMin2[1])\n"
" return 0;\n"
" if (aabbMin1[2] > aabbMax2[2])\n"
" return 0;\n"
" if (aabbMax1[2] < aabbMin2[2])\n"
" return 0;\n"
" return 1;\n"
" //overlap = ((aabbMin1[0] > aabbMax2[0]) || (aabbMax1[0] < aabbMin2[0])) ? 0 : overlap;\n"
" //overlap = ((aabbMin1[2] > aabbMax2[2]) || (aabbMax1[2] < aabbMin2[2])) ? 0 : overlap;\n"
" //overlap = ((aabbMin1[1] > aabbMax2[1]) || (aabbMax1[1] < aabbMin2[1])) ? 0 : overlap;\n"
" //return overlap;\n"
"}\n"
"\n"
"\n"
@@ -178,88 +195,78 @@ static const char* bvhTraversalKernelCL= \
" int numPairs,\n"
" int maxNumConcavePairsCapacity)\n"
"{\n"
"\n"
" int i = get_global_id(0);\n"
" int id = get_global_id(0);\n"
" if (id>=numPairs)\n"
" return;\n"
" \n"
" if (i<numPairs)\n"
" int bodyIndexA = pairs[id].x;\n"
" int bodyIndexB = pairs[id].y;\n"
" int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
" int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
" \n"
" //once the broadphase avoids static-static pairs, we can remove this test\n"
" if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))\n"
" {\n"
" return;\n"
" }\n"
" \n"
" if (collidables[collidableIndexA].m_shapeType!=SHAPE_CONCAVE_TRIMESH)\n"
" return;\n"
"\n"
" if (collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL)\n"
" return;\n"
"\n"
" unsigned short int quantizedQueryAabbMin[3];\n"
" unsigned short int quantizedQueryAabbMax[3];\n"
" quantizeWithClamp(quantizedQueryAabbMin,aabbs[bodyIndexB].m_min,false,bvhAabbMin, bvhAabbMax,bvhQuantization);\n"
" quantizeWithClamp(quantizedQueryAabbMax,aabbs[bodyIndexB].m_max,true ,bvhAabbMin, bvhAabbMax,bvhQuantization);\n"
" \n"
" int bodyIndexA = pairs[i].x;\n"
" int bodyIndexB = pairs[i].y;\n"
"\n"
" int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
" int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
" \n"
" int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;\n"
" int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;\n"
" \n"
" \n"
" //once the broadphase avoids static-static pairs, we can remove this test\n"
" if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))\n"
" for (int i=0;i<numSubtreeHeaders;i++)\n"
" {\n"
" btBvhSubtreeInfo subtree = subtreeHeaders[i];\n"
" \n"
" int overlap = testQuantizedAabbAgainstQuantizedAabb(quantizedQueryAabbMin,quantizedQueryAabbMax,subtree.m_quantizedAabbMin,subtree.m_quantizedAabbMax);\n"
" if (overlap != 0)\n"
" {\n"
" return;\n"
" }\n"
" \n"
" if ((collidables[collidableIndexA].m_shapeType==SHAPE_CONCAVE_TRIMESH))// && (collidables[collidableIndexB].m_shapeType==SHAPE_CONVEX_HULL))\n"
" {\n"
"\n"
" \n"
" unsigned short int quantizedQueryAabbMin[3];\n"
" unsigned short int quantizedQueryAabbMax[3];\n"
" quantizeWithClamp(quantizedQueryAabbMin,aabbs[bodyIndexB].m_min,false,bvhAabbMin, bvhAabbMax,bvhQuantization);\n"
" quantizeWithClamp(quantizedQueryAabbMax,aabbs[bodyIndexB].m_max,true ,bvhAabbMin, bvhAabbMax,bvhQuantization);\n"
"\n"
"\n"
" int i;\n"
" for (i=0;i<numSubtreeHeaders;i++)\n"
" int startNodeIndex = subtree.m_rootNodeIndex;\n"
" int endNodeIndex = subtree.m_rootNodeIndex+subtree.m_subtreeSize;\n"
" int curIndex = startNodeIndex;\n"
" int escapeIndex;\n"
" int isLeafNode;\n"
" int aabbOverlap;\n"
" while (curIndex < endNodeIndex)\n"
" {\n"
" const __global btBvhSubtreeInfo* subtree = &subtreeHeaders[i];\n"
" //PCK: unsigned instead of bool\n"
" unsigned overlap = testQuantizedAabbAgainstQuantizedAabb(quantizedQueryAabbMin,quantizedQueryAabbMax,subtree->m_quantizedAabbMin,subtree->m_quantizedAabbMax);\n"
" if (overlap != 0)\n"
" btQuantizedBvhNode rootNode = quantizedNodes[curIndex];\n"
" aabbOverlap = testQuantizedAabbAgainstQuantizedAabb(quantizedQueryAabbMin,quantizedQueryAabbMax,rootNode.m_quantizedAabbMin,rootNode.m_quantizedAabbMax);\n"
" isLeafNode = isLeaf(&rootNode);\n"
" if (aabbOverlap)\n"
" {\n"
" int startNodeIndex = subtree->m_rootNodeIndex;\n"
" int endNodeIndex = subtree->m_rootNodeIndex+subtree->m_subtreeSize;\n"
"\n"
" int curIndex = startNodeIndex;\n"
" int subTreeSize = endNodeIndex - startNodeIndex;\n"
" __global const btQuantizedBvhNode* rootNode = &quantizedNodes[startNodeIndex];\n"
" int escapeIndex;\n"
" bool isLeafNode;\n"
" unsigned aabbOverlap;\n"
" while (curIndex < endNodeIndex)\n"
" if (isLeafNode)\n"
" {\n"
" aabbOverlap = testQuantizedAabbAgainstQuantizedAabb(quantizedQueryAabbMin,quantizedQueryAabbMax,rootNode->m_quantizedAabbMin,rootNode->m_quantizedAabbMax);\n"
" isLeafNode = isLeaf(rootNode);\n"
" if (isLeafNode && aabbOverlap)\n"
" int triangleIndex = getTriangleIndex(&rootNode);\n"
" \n"
" int pairIdx = atomic_inc(numConcavePairsOut);\n"
" if (pairIdx<maxNumConcavePairsCapacity)\n"
" {\n"
" //do your thing! nodeCallback->processNode(rootNode->getPartId(),rootNode->getTriangleIndex());\n"
" int triangleIndex = getTriangleIndex(rootNode);\n"
" int pairIdx = atomic_inc(numConcavePairsOut);\n"
" if (pairIdx<maxNumConcavePairsCapacity)\n"
" {\n"
" //int4 newPair;\n"
" concavePairsOut[pairIdx].x = bodyIndexA;\n"
" concavePairsOut[pairIdx].y = bodyIndexB;\n"
" concavePairsOut[pairIdx].z = triangleIndex;\n"
" concavePairsOut[pairIdx].w = 3;\n"
" }\n"
" } \n"
" if ((aabbOverlap != 0) || isLeafNode)\n"
" {\n"
" rootNode++;\n"
" curIndex++;\n"
" } else\n"
" {\n"
" escapeIndex = getEscapeIndex(rootNode);\n"
" rootNode += escapeIndex;\n"
" curIndex += escapeIndex;\n"
" int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,3);\n"
" concavePairsOut[pairIdx] = newPair;\n"
" }\n"
" } \n"
" curIndex++;\n"
" } else\n"
" {\n"
" if (isLeafNode)\n"
" {\n"
" curIndex++;\n"
" } else\n"
" {\n"
" escapeIndex = getEscapeIndex(&rootNode);\n"
" curIndex += escapeIndex;\n"
" }\n"
" }\n"
" }\n"
" }//SHAPE_CONCAVE_TRIMESH\n"
" }//i<numpairs\n"
" }\n"
" }\n"
"\n"
"}\n"
;