2017-08-01 14:30:58 +02:00
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
2019-01-03 14:26:51 +01:00
static const char * bvhTraversalKernelCL =
" //keep this enum in sync with the CPU version (in btCollidable.h) \n "
" //written by Erwin Coumans \n "
" #define SHAPE_CONVEX_HULL 3 \n "
" #define SHAPE_CONCAVE_TRIMESH 5 \n "
" #define TRIANGLE_NUM_CONVEX_FACES 5 \n "
" #define SHAPE_COMPOUND_OF_CONVEX_HULLS 6 \n "
" #define SHAPE_SPHERE 7 \n "
" typedef unsigned int u32; \n "
" #define MAX_NUM_PARTS_IN_BITS 10 \n "
" ///btQuantizedBvhNode is a compressed aabb node, 16 bytes. \n "
" ///Node can be used for leafnode or internal node. Leafnodes can point to 32-bit triangle index (non-negative range). \n "
" typedef struct \n "
" { \n "
" //12 bytes \n "
" unsigned short int m_quantizedAabbMin[3]; \n "
" unsigned short int m_quantizedAabbMax[3]; \n "
" //4 bytes \n "
" int m_escapeIndexOrTriangleIndex; \n "
" } btQuantizedBvhNode; \n "
" typedef struct \n "
" { \n "
" float4 m_aabbMin; \n "
" float4 m_aabbMax; \n "
" float4 m_quantization; \n "
" int m_numNodes; \n "
" int m_numSubTrees; \n "
" int m_nodeOffset; \n "
" int m_subTreeOffset; \n "
" } b3BvhInfo; \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 "
" // Get only the lower bits where the triangle index is stored \n "
" return (rootNode->m_escapeIndexOrTriangleIndex&~(y)); \n "
" } \n "
" int isLeaf(const btQuantizedBvhNode* rootNode) \n "
" { \n "
" //skipindex is negative (internal node), triangleindex >=0 (leafnode) \n "
" return (rootNode->m_escapeIndexOrTriangleIndex >= 0)? 1 : 0; \n "
" } \n "
" \n "
" int getEscapeIndex(const btQuantizedBvhNode* rootNode) \n "
" { \n "
" return -rootNode->m_escapeIndexOrTriangleIndex; \n "
" } \n "
" typedef struct \n "
" { \n "
" //12 bytes \n "
" unsigned short int m_quantizedAabbMin[3]; \n "
" unsigned short int m_quantizedAabbMax[3]; \n "
" //4 bytes, points to the root of the subtree \n "
" int m_rootNodeIndex; \n "
" //4 bytes \n "
" int m_subtreeSize; \n "
" int m_padding[3]; \n "
" } btBvhSubtreeInfo; \n "
" ///keep this in sync with btCollidable.h \n "
" typedef struct \n "
" { \n "
" int m_numChildShapes; \n "
" int blaat2; \n "
" int m_shapeType; \n "
" int m_shapeIndex; \n "
" \n "
" } btCollidableGpu; \n "
" typedef struct \n "
" { \n "
" float4 m_childPosition; \n "
" float4 m_childOrientation; \n "
" int m_shapeIndex; \n "
" int m_unused0; \n "
" int m_unused1; \n "
" int m_unused2; \n "
" } btGpuChildShape; \n "
" typedef struct \n "
" { \n "
" float4 m_pos; \n "
" float4 m_quat; \n "
" float4 m_linVel; \n "
" float4 m_angVel; \n "
" u32 m_collidableIdx; \n "
" float m_invMass; \n "
" float m_restituitionCoeff; \n "
" float m_frictionCoeff; \n "
" } BodyData; \n "
" typedef struct \n "
" { \n "
" union \n "
" { \n "
" float4 m_min; \n "
" float m_minElems[4]; \n "
" int m_minIndices[4]; \n "
" }; \n "
" union \n "
" { \n "
" float4 m_max; \n "
" float m_maxElems[4]; \n "
" int m_maxIndices[4]; \n "
" }; \n "
" } btAabbCL; \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 "
" //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 "
" void quantizeWithClamp(unsigned short* out, float4 point2,int isMax, float4 bvhAabbMin, float4 bvhAabbMax, float4 bvhQuantization) \n "
" { \n "
" float4 clampedPoint = max(point2,bvhAabbMin); \n "
" clampedPoint = min (clampedPoint, bvhAabbMax); \n "
" float4 v = (clampedPoint - bvhAabbMin) * bvhQuantization; \n "
" if (isMax) \n "
" { \n "
" out[0] = (unsigned short) (((unsigned short)(v.x+1.f) | 1)); \n "
" out[1] = (unsigned short) (((unsigned short)(v.y+1.f) | 1)); \n "
" out[2] = (unsigned short) (((unsigned short)(v.z+1.f) | 1)); \n "
" } else \n "
" { \n "
" out[0] = (unsigned short) (((unsigned short)(v.x) & 0xfffe)); \n "
" out[1] = (unsigned short) (((unsigned short)(v.y) & 0xfffe)); \n "
" out[2] = (unsigned short) (((unsigned short)(v.z) & 0xfffe)); \n "
" } \n "
" } \n "
" // work-in-progress \n "
" __kernel void bvhTraversalKernel( __global const int4* pairs, \n "
" __global const BodyData* rigidBodies, \n "
" __global const btCollidableGpu* collidables, \n "
" __global btAabbCL* aabbs, \n "
" __global int4* concavePairsOut, \n "
" __global volatile int* numConcavePairsOut, \n "
" __global const btBvhSubtreeInfo* subtreeHeadersRoot, \n "
" __global const btQuantizedBvhNode* quantizedNodesRoot, \n "
" __global const b3BvhInfo* bvhInfos, \n "
" int numPairs, \n "
" int maxNumConcavePairsCapacity) \n "
" { \n "
" int id = get_global_id(0); \n "
" if (id>=numPairs) \n "
" return; \n "
" \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 "
" int shapeTypeB = collidables[collidableIndexB].m_shapeType; \n "
" \n "
" if (shapeTypeB!=SHAPE_CONVEX_HULL && \n "
" shapeTypeB!=SHAPE_SPHERE && \n "
" shapeTypeB!=SHAPE_COMPOUND_OF_CONVEX_HULLS \n "
" ) \n "
" return; \n "
" b3BvhInfo bvhInfo = bvhInfos[collidables[collidableIndexA].m_numChildShapes]; \n "
" float4 bvhAabbMin = bvhInfo.m_aabbMin; \n "
" float4 bvhAabbMax = bvhInfo.m_aabbMax; \n "
" float4 bvhQuantization = bvhInfo.m_quantization; \n "
" int numSubtreeHeaders = bvhInfo.m_numSubTrees; \n "
" __global const btBvhSubtreeInfo* subtreeHeaders = &subtreeHeadersRoot[bvhInfo.m_subTreeOffset]; \n "
" __global const btQuantizedBvhNode* quantizedNodes = &quantizedNodesRoot[bvhInfo.m_nodeOffset]; \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 "
" 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 "
" 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 "
" btQuantizedBvhNode rootNode = quantizedNodes[curIndex]; \n "
" aabbOverlap = testQuantizedAabbAgainstQuantizedAabb(quantizedQueryAabbMin,quantizedQueryAabbMax,rootNode.m_quantizedAabbMin,rootNode.m_quantizedAabbMax); \n "
" isLeafNode = isLeaf(&rootNode); \n "
" if (aabbOverlap) \n "
" { \n "
" if (isLeafNode) \n "
" { \n "
" int triangleIndex = getTriangleIndex(&rootNode); \n "
" if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS) \n "
" { \n "
" int numChildrenB = collidables[collidableIndexB].m_numChildShapes; \n "
" int pairIdx = atomic_add(numConcavePairsOut,numChildrenB); \n "
" for (int b=0;b<numChildrenB;b++) \n "
" { \n "
" if ((pairIdx+b)<maxNumConcavePairsCapacity) \n "
" { \n "
" int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+b; \n "
" int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,childShapeIndexB); \n "
" concavePairsOut[pairIdx+b] = newPair; \n "
" } \n "
" } \n "
" } else \n "
" { \n "
" int pairIdx = atomic_inc(numConcavePairsOut); \n "
" if (pairIdx<maxNumConcavePairsCapacity) \n "
" { \n "
" int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,0); \n "
" concavePairsOut[pairIdx] = newPair; \n "
" } \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 "
" } \n "
" } \n "
" } \n " ;
