dali-physics/third-party/bullet3/src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.h

   1 //this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
   2 static const char* bvhTraversalKernelCL =
   3         "//keep this enum in sync with the CPU version (in btCollidable.h)\n"
   4         "//written by Erwin Coumans\n"
   5         "#define SHAPE_CONVEX_HULL 3\n"
   6         "#define SHAPE_CONCAVE_TRIMESH 5\n"
   7         "#define TRIANGLE_NUM_CONVEX_FACES 5\n"
   8         "#define SHAPE_COMPOUND_OF_CONVEX_HULLS 6\n"
   9         "#define SHAPE_SPHERE 7\n"
  10         "typedef unsigned int u32;\n"
  11         "#define MAX_NUM_PARTS_IN_BITS 10\n"
  12         "///btQuantizedBvhNode is a compressed aabb node, 16 bytes.\n"
  13         "///Node can be used for leafnode or internal node. Leafnodes can point to 32-bit triangle index (non-negative range).\n"
  14         "typedef struct\n"
  15         "{\n"
  16         "       //12 bytes\n"
  17         "       unsigned short int      m_quantizedAabbMin[3];\n"
  18         "       unsigned short int      m_quantizedAabbMax[3];\n"
  19         "       //4 bytes\n"
  20         "       int     m_escapeIndexOrTriangleIndex;\n"
  21         "} btQuantizedBvhNode;\n"
  22         "typedef struct\n"
  23         "{\n"
  24         "       float4          m_aabbMin;\n"
  25         "       float4          m_aabbMax;\n"
  26         "       float4          m_quantization;\n"
  27         "       int                     m_numNodes;\n"
  28         "       int                     m_numSubTrees;\n"
  29         "       int                     m_nodeOffset;\n"
  30         "       int                     m_subTreeOffset;\n"
  31         "} b3BvhInfo;\n"
  32         "int    getTriangleIndex(const btQuantizedBvhNode* rootNode)\n"
  33         "{\n"
  34         "       unsigned int x=0;\n"
  35         "       unsigned int y = (~(x&0))<<(31-MAX_NUM_PARTS_IN_BITS);\n"
  36         "       // Get only the lower bits where the triangle index is stored\n"
  37         "       return (rootNode->m_escapeIndexOrTriangleIndex&~(y));\n"
  38         "}\n"
  39         "int isLeaf(const btQuantizedBvhNode* rootNode)\n"
  40         "{\n"
  41         "       //skipindex is negative (internal node), triangleindex >=0 (leafnode)\n"
  42         "       return (rootNode->m_escapeIndexOrTriangleIndex >= 0)? 1 : 0;\n"
  43         "}\n"
  44         "       \n"
  45         "int getEscapeIndex(const btQuantizedBvhNode* rootNode)\n"
  46         "{\n"
  47         "       return -rootNode->m_escapeIndexOrTriangleIndex;\n"
  48         "}\n"
  49         "typedef struct\n"
  50         "{\n"
  51         "       //12 bytes\n"
  52         "       unsigned short int      m_quantizedAabbMin[3];\n"
  53         "       unsigned short int      m_quantizedAabbMax[3];\n"
  54         "       //4 bytes, points to the root of the subtree\n"
  55         "       int                     m_rootNodeIndex;\n"
  56         "       //4 bytes\n"
  57         "       int                     m_subtreeSize;\n"
  58         "       int                     m_padding[3];\n"
  59         "} btBvhSubtreeInfo;\n"
  60         "///keep this in sync with btCollidable.h\n"
  61         "typedef struct\n"
  62         "{\n"
  63         "       int m_numChildShapes;\n"
  64         "       int blaat2;\n"
  65         "       int m_shapeType;\n"
  66         "       int m_shapeIndex;\n"
  67         "       \n"
  68         "} btCollidableGpu;\n"
  69         "typedef struct\n"
  70         "{\n"
  71         "       float4  m_childPosition;\n"
  72         "       float4  m_childOrientation;\n"
  73         "       int m_shapeIndex;\n"
  74         "       int m_unused0;\n"
  75         "       int m_unused1;\n"
  76         "       int m_unused2;\n"
  77         "} btGpuChildShape;\n"
  78         "typedef struct\n"
  79         "{\n"
  80         "       float4 m_pos;\n"
  81         "       float4 m_quat;\n"
  82         "       float4 m_linVel;\n"
  83         "       float4 m_angVel;\n"
  84         "       u32 m_collidableIdx;\n"
  85         "       float m_invMass;\n"
  86         "       float m_restituitionCoeff;\n"
  87         "       float m_frictionCoeff;\n"
  88         "} BodyData;\n"
  89         "typedef struct \n"
  90         "{\n"
  91         "       union\n"
  92         "       {\n"
  93         "               float4  m_min;\n"
  94         "               float   m_minElems[4];\n"
  95         "               int                     m_minIndices[4];\n"
  96         "       };\n"
  97         "       union\n"
  98         "       {\n"
  99         "               float4  m_max;\n"
 100         "               float   m_maxElems[4];\n"
 101         "               int                     m_maxIndices[4];\n"
 102         "       };\n"
 103         "} btAabbCL;\n"
 104         "int testQuantizedAabbAgainstQuantizedAabb(\n"
 105         "                                                               const unsigned short int* aabbMin1,\n"
 106         "                                                               const unsigned short int* aabbMax1,\n"
 107         "                                                               const unsigned short int* aabbMin2,\n"
 108         "                                                               const unsigned short int* aabbMax2)\n"
 109         "{\n"
 110         "       //int overlap = 1;\n"
 111         "       if (aabbMin1[0] > aabbMax2[0])\n"
 112         "               return 0;\n"
 113         "       if (aabbMax1[0] < aabbMin2[0])\n"
 114         "               return 0;\n"
 115         "       if (aabbMin1[1] > aabbMax2[1])\n"
 116         "               return 0;\n"
 117         "       if (aabbMax1[1] < aabbMin2[1])\n"
 118         "               return 0;\n"
 119         "       if (aabbMin1[2] > aabbMax2[2])\n"
 120         "               return 0;\n"
 121         "       if (aabbMax1[2] < aabbMin2[2])\n"
 122         "               return 0;\n"
 123         "       return 1;\n"
 124         "       //overlap = ((aabbMin1[0] > aabbMax2[0]) || (aabbMax1[0] < aabbMin2[0])) ? 0 : overlap;\n"
 125         "       //overlap = ((aabbMin1[2] > aabbMax2[2]) || (aabbMax1[2] < aabbMin2[2])) ? 0 : overlap;\n"
 126         "       //overlap = ((aabbMin1[1] > aabbMax2[1]) || (aabbMax1[1] < aabbMin2[1])) ? 0 : overlap;\n"
 127         "       //return overlap;\n"
 128         "}\n"
 129         "void quantizeWithClamp(unsigned short* out, float4 point2,int isMax, float4 bvhAabbMin, float4 bvhAabbMax, float4 bvhQuantization)\n"
 130         "{\n"
 131         "       float4 clampedPoint = max(point2,bvhAabbMin);\n"
 132         "       clampedPoint = min (clampedPoint, bvhAabbMax);\n"
 133         "       float4 v = (clampedPoint - bvhAabbMin) * bvhQuantization;\n"
 134         "       if (isMax)\n"
 135         "       {\n"
 136         "               out[0] = (unsigned short) (((unsigned short)(v.x+1.f) | 1));\n"
 137         "               out[1] = (unsigned short) (((unsigned short)(v.y+1.f) | 1));\n"
 138         "               out[2] = (unsigned short) (((unsigned short)(v.z+1.f) | 1));\n"
 139         "       } else\n"
 140         "       {\n"
 141         "               out[0] = (unsigned short) (((unsigned short)(v.x) & 0xfffe));\n"
 142         "               out[1] = (unsigned short) (((unsigned short)(v.y) & 0xfffe));\n"
 143         "               out[2] = (unsigned short) (((unsigned short)(v.z) & 0xfffe));\n"
 144         "       }\n"
 145         "}\n"
 146         "// work-in-progress\n"
 147         "__kernel void   bvhTraversalKernel( __global const int4* pairs, \n"
 148         "                                                                       __global const BodyData* rigidBodies, \n"
 149         "                                                                       __global const btCollidableGpu* collidables,\n"
 150         "                                                                       __global btAabbCL* aabbs,\n"
 151         "                                                                       __global int4* concavePairsOut,\n"
 152         "                                                                       __global volatile int* numConcavePairsOut,\n"
 153         "                                                                       __global const btBvhSubtreeInfo* subtreeHeadersRoot,\n"
 154         "                                                                       __global const btQuantizedBvhNode* quantizedNodesRoot,\n"
 155         "                                                                       __global const b3BvhInfo* bvhInfos,\n"
 156         "                                                                       int numPairs,\n"
 157         "                                                                       int maxNumConcavePairsCapacity)\n"
 158         "{\n"
 159         "       int id = get_global_id(0);\n"
 160         "       if (id>=numPairs)\n"
 161         "               return;\n"
 162         "       \n"
 163         "       int bodyIndexA = pairs[id].x;\n"
 164         "       int bodyIndexB = pairs[id].y;\n"
 165         "       int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
 166         "       int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
 167         "       \n"
 168         "       //once the broadphase avoids static-static pairs, we can remove this test\n"
 169         "       if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))\n"
 170         "       {\n"
 171         "               return;\n"
 172         "       }\n"
 173         "               \n"
 174         "       if (collidables[collidableIndexA].m_shapeType!=SHAPE_CONCAVE_TRIMESH)\n"
 175         "               return;\n"
 176         "       int shapeTypeB = collidables[collidableIndexB].m_shapeType;\n"
 177         "               \n"
 178         "       if (shapeTypeB!=SHAPE_CONVEX_HULL &&\n"
 179         "               shapeTypeB!=SHAPE_SPHERE        &&\n"
 180         "               shapeTypeB!=SHAPE_COMPOUND_OF_CONVEX_HULLS\n"
 181         "               )\n"
 182         "               return;\n"
 183         "       b3BvhInfo bvhInfo = bvhInfos[collidables[collidableIndexA].m_numChildShapes];\n"
 184         "       float4 bvhAabbMin = bvhInfo.m_aabbMin;\n"
 185         "       float4 bvhAabbMax = bvhInfo.m_aabbMax;\n"
 186         "       float4 bvhQuantization = bvhInfo.m_quantization;\n"
 187         "       int numSubtreeHeaders = bvhInfo.m_numSubTrees;\n"
 188         "       __global const btBvhSubtreeInfo* subtreeHeaders = &subtreeHeadersRoot[bvhInfo.m_subTreeOffset];\n"
 189         "       __global const btQuantizedBvhNode* quantizedNodes = &quantizedNodesRoot[bvhInfo.m_nodeOffset];\n"
 190         "       \n"
 191         "       unsigned short int quantizedQueryAabbMin[3];\n"
 192         "       unsigned short int quantizedQueryAabbMax[3];\n"
 193         "       quantizeWithClamp(quantizedQueryAabbMin,aabbs[bodyIndexB].m_min,false,bvhAabbMin, bvhAabbMax,bvhQuantization);\n"
 194         "       quantizeWithClamp(quantizedQueryAabbMax,aabbs[bodyIndexB].m_max,true ,bvhAabbMin, bvhAabbMax,bvhQuantization);\n"
 195         "       \n"
 196         "       for (int i=0;i<numSubtreeHeaders;i++)\n"
 197         "       {\n"
 198         "               btBvhSubtreeInfo subtree = subtreeHeaders[i];\n"
 199         "                               \n"
 200         "               int overlap = testQuantizedAabbAgainstQuantizedAabb(quantizedQueryAabbMin,quantizedQueryAabbMax,subtree.m_quantizedAabbMin,subtree.m_quantizedAabbMax);\n"
 201         "               if (overlap != 0)\n"
 202         "               {\n"
 203         "                       int startNodeIndex = subtree.m_rootNodeIndex;\n"
 204         "                       int endNodeIndex = subtree.m_rootNodeIndex+subtree.m_subtreeSize;\n"
 205         "                       int curIndex = startNodeIndex;\n"
 206         "                       int escapeIndex;\n"
 207         "                       int isLeafNode;\n"
 208         "                       int aabbOverlap;\n"
 209         "                       while (curIndex < endNodeIndex)\n"
 210         "                       {\n"
 211         "                               btQuantizedBvhNode rootNode = quantizedNodes[curIndex];\n"
 212         "                               aabbOverlap = testQuantizedAabbAgainstQuantizedAabb(quantizedQueryAabbMin,quantizedQueryAabbMax,rootNode.m_quantizedAabbMin,rootNode.m_quantizedAabbMax);\n"
 213         "                               isLeafNode = isLeaf(&rootNode);\n"
 214         "                               if (aabbOverlap)\n"
 215         "                               {\n"
 216         "                                       if (isLeafNode)\n"
 217         "                                       {\n"
 218         "                                               int triangleIndex = getTriangleIndex(&rootNode);\n"
 219         "                                               if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
 220         "                                               {\n"
 221         "                                                               int numChildrenB = collidables[collidableIndexB].m_numChildShapes;\n"
 222         "                                                               int pairIdx = atomic_add(numConcavePairsOut,numChildrenB);\n"
 223         "                                                               for (int b=0;b<numChildrenB;b++)\n"
 224         "                                                               {\n"
 225         "                                                                       if ((pairIdx+b)<maxNumConcavePairsCapacity)\n"
 226         "                                                                       {\n"
 227         "                                                                               int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+b;\n"
 228         "                                                                               int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,childShapeIndexB);\n"
 229         "                                                                               concavePairsOut[pairIdx+b] = newPair;\n"
 230         "                                                                       }\n"
 231         "                                                               }\n"
 232         "                                               } else\n"
 233         "                                               {\n"
 234         "                                                       int pairIdx = atomic_inc(numConcavePairsOut);\n"
 235         "                                                       if (pairIdx<maxNumConcavePairsCapacity)\n"
 236         "                                                       {\n"
 237         "                                                               int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,0);\n"
 238         "                                                               concavePairsOut[pairIdx] = newPair;\n"
 239         "                                                       }\n"
 240         "                                               }\n"
 241         "                                       } \n"
 242         "                                       curIndex++;\n"
 243         "                               } else\n"
 244         "                               {\n"
 245         "                                       if (isLeafNode)\n"
 246         "                                       {\n"
 247         "                                               curIndex++;\n"
 248         "                                       } else\n"
 249         "                                       {\n"
 250         "                                               escapeIndex = getEscapeIndex(&rootNode);\n"
 251         "                                               curIndex += escapeIndex;\n"
 252         "                                       }\n"
 253         "                               }\n"
 254         "                       }\n"
 255         "               }\n"
 256         "       }\n"
 257         "}\n";