dali-physics/third-party/bullet3/src/Bullet3OpenCL/BroadphaseCollision/kernels/gridBroadphaseKernels.h

   1 //this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
   2 static const char* gridBroadphaseCL =
   3         "int getPosHash(int4 gridPos, __global float4* pParams)\n"
   4         "{\n"
   5         "       int4 gridDim = *((__global int4*)(pParams + 1));\n"
   6         "       gridPos.x &= gridDim.x - 1;\n"
   7         "       gridPos.y &= gridDim.y - 1;\n"
   8         "       gridPos.z &= gridDim.z - 1;\n"
   9         "       int hash = gridPos.z * gridDim.y * gridDim.x + gridPos.y * gridDim.x + gridPos.x;\n"
  10         "       return hash;\n"
  11         "} \n"
  12         "int4 getGridPos(float4 worldPos, __global float4* pParams)\n"
  13         "{\n"
  14         "    int4 gridPos;\n"
  15         "       int4 gridDim = *((__global int4*)(pParams + 1));\n"
  16         "    gridPos.x = (int)floor(worldPos.x * pParams[0].x) & (gridDim.x - 1);\n"
  17         "    gridPos.y = (int)floor(worldPos.y * pParams[0].y) & (gridDim.y - 1);\n"
  18         "    gridPos.z = (int)floor(worldPos.z * pParams[0].z) & (gridDim.z - 1);\n"
  19         "    return gridPos;\n"
  20         "}\n"
  21         "// calculate grid hash value for each body using its AABB\n"
  22         "__kernel void kCalcHashAABB(int numObjects, __global float4* allpAABB, __global const int* smallAabbMapping, __global int2* pHash, __global float4* pParams )\n"
  23         "{\n"
  24         "    int index = get_global_id(0);\n"
  25         "    if(index >= numObjects)\n"
  26         "       {\n"
  27         "               return;\n"
  28         "       }\n"
  29         "       float4 bbMin = allpAABB[smallAabbMapping[index]*2];\n"
  30         "       float4 bbMax = allpAABB[smallAabbMapping[index]*2 + 1];\n"
  31         "       float4 pos;\n"
  32         "       pos.x = (bbMin.x + bbMax.x) * 0.5f;\n"
  33         "       pos.y = (bbMin.y + bbMax.y) * 0.5f;\n"
  34         "       pos.z = (bbMin.z + bbMax.z) * 0.5f;\n"
  35         "       pos.w = 0.f;\n"
  36         "    // get address in grid\n"
  37         "    int4 gridPos = getGridPos(pos, pParams);\n"
  38         "    int gridHash = getPosHash(gridPos, pParams);\n"
  39         "    // store grid hash and body index\n"
  40         "    int2 hashVal;\n"
  41         "    hashVal.x = gridHash;\n"
  42         "    hashVal.y = index;\n"
  43         "    pHash[index] = hashVal;\n"
  44         "}\n"
  45         "__kernel void kClearCellStart( int numCells, \n"
  46         "                                                               __global int* pCellStart )\n"
  47         "{\n"
  48         "    int index = get_global_id(0);\n"
  49         "    if(index >= numCells)\n"
  50         "       {\n"
  51         "               return;\n"
  52         "       }\n"
  53         "       pCellStart[index] = -1;\n"
  54         "}\n"
  55         "__kernel void kFindCellStart(int numObjects, __global int2* pHash, __global int* cellStart )\n"
  56         "{\n"
  57         "       __local int sharedHash[513];\n"
  58         "    int index = get_global_id(0);\n"
  59         "       int2 sortedData;\n"
  60         "    if(index < numObjects)\n"
  61         "       {\n"
  62         "               sortedData = pHash[index];\n"
  63         "               // Load hash data into shared memory so that we can look \n"
  64         "               // at neighboring body's hash value without loading\n"
  65         "               // two hash values per thread\n"
  66         "               sharedHash[get_local_id(0) + 1] = sortedData.x;\n"
  67         "               if((index > 0) && (get_local_id(0) == 0))\n"
  68         "               {\n"
  69         "                       // first thread in block must load neighbor body hash\n"
  70         "                       sharedHash[0] = pHash[index-1].x;\n"
  71         "               }\n"
  72         "       }\n"
  73         "    barrier(CLK_LOCAL_MEM_FENCE);\n"
  74         "    if(index < numObjects)\n"
  75         "       {\n"
  76         "               if((index == 0) || (sortedData.x != sharedHash[get_local_id(0)]))\n"
  77         "               {\n"
  78         "                       cellStart[sortedData.x] = index;\n"
  79         "               }\n"
  80         "       }\n"
  81         "}\n"
  82         "int testAABBOverlap(float4 min0, float4 max0, float4 min1, float4 max1)\n"
  83         "{\n"
  84         "       return  (min0.x <= max1.x)&& (min1.x <= max0.x) && \n"
  85         "                       (min0.y <= max1.y)&& (min1.y <= max0.y) && \n"
  86         "                       (min0.z <= max1.z)&& (min1.z <= max0.z); \n"
  87         "}\n"
  88         "//search for AABB 'index' against other AABBs' in this cell\n"
  89         "void findPairsInCell(  int numObjects,\n"
  90         "                                               int4    gridPos,\n"
  91         "                                               int    index,\n"
  92         "                                               __global int2*  pHash,\n"
  93         "                                               __global int*   pCellStart,\n"
  94         "                                               __global float4* allpAABB, \n"
  95         "                                               __global const int* smallAabbMapping,\n"
  96         "                                               __global float4* pParams,\n"
  97         "                                                       volatile  __global int* pairCount,\n"
  98         "                                               __global int4*   pPairBuff2,\n"
  99         "                                               int maxPairs\n"
 100         "                                               )\n"
 101         "{\n"
 102         "       int4 pGridDim = *((__global int4*)(pParams + 1));\n"
 103         "       int maxBodiesPerCell = pGridDim.w;\n"
 104         "    int gridHash = getPosHash(gridPos, pParams);\n"
 105         "    // get start of bucket for this cell\n"
 106         "    int bucketStart = pCellStart[gridHash];\n"
 107         "    if (bucketStart == -1)\n"
 108         "       {\n"
 109         "        return;   // cell empty\n"
 110         "       }\n"
 111         "       // iterate over bodies in this cell\n"
 112         "    int2 sortedData = pHash[index];\n"
 113         "       int unsorted_indx = sortedData.y;\n"
 114         "    float4 min0 = allpAABB[smallAabbMapping[unsorted_indx]*2 + 0]; \n"
 115         "       float4 max0 = allpAABB[smallAabbMapping[unsorted_indx]*2 + 1];\n"
 116         "       int handleIndex =  as_int(min0.w);\n"
 117         "       \n"
 118         "       int bucketEnd = bucketStart + maxBodiesPerCell;\n"
 119         "       bucketEnd = (bucketEnd > numObjects) ? numObjects : bucketEnd;\n"
 120         "       for(int index2 = bucketStart; index2 < bucketEnd; index2++) \n"
 121         "       {\n"
 122         "        int2 cellData = pHash[index2];\n"
 123         "        if (cellData.x != gridHash)\n"
 124         "        {\n"
 125         "                       break;   // no longer in same bucket\n"
 126         "               }\n"
 127         "               int unsorted_indx2 = cellData.y;\n"
 128         "        //if (unsorted_indx2 < unsorted_indx) // check not colliding with self\n"
 129         "               if (unsorted_indx2 != unsorted_indx) // check not colliding with self\n"
 130         "        {   \n"
 131         "                       float4 min1 = allpAABB[smallAabbMapping[unsorted_indx2]*2 + 0];\n"
 132         "                       float4 max1 = allpAABB[smallAabbMapping[unsorted_indx2]*2 + 1];\n"
 133         "                       if(testAABBOverlap(min0, max0, min1, max1))\n"
 134         "                       {\n"
 135         "                               if (pairCount)\n"
 136         "                               {\n"
 137         "                                       int handleIndex2 = as_int(min1.w);\n"
 138         "                                       if (handleIndex<handleIndex2)\n"
 139         "                                       {\n"
 140         "                                               int curPair = atomic_add(pairCount,1);\n"
 141         "                                               if (curPair<maxPairs)\n"
 142         "                                               {\n"
 143         "                                                       int4 newpair;\n"
 144         "                                                       newpair.x = handleIndex;\n"
 145         "                                                       newpair.y = handleIndex2;\n"
 146         "                                                       newpair.z = -1;\n"
 147         "                                                       newpair.w = -1;\n"
 148         "                                                       pPairBuff2[curPair] = newpair;\n"
 149         "                                               }\n"
 150         "                                       }\n"
 151         "                               \n"
 152         "                               }\n"
 153         "                       }\n"
 154         "               }\n"
 155         "       }\n"
 156         "}\n"
 157         "__kernel void kFindOverlappingPairs(   int numObjects,\n"
 158         "                                                                               __global float4* allpAABB, \n"
 159         "                                                                               __global const int* smallAabbMapping,\n"
 160         "                                                                               __global int2* pHash, \n"
 161         "                                                                               __global int* pCellStart, \n"
 162         "                                                                               __global float4* pParams ,\n"
 163         "                                                                               volatile  __global int* pairCount,\n"
 164         "                                                                               __global int4*   pPairBuff2,\n"
 165         "                                                                               int maxPairs\n"
 166         "                                                                               )\n"
 167         "{\n"
 168         "    int index = get_global_id(0);\n"
 169         "    if(index >= numObjects)\n"
 170         "       {\n"
 171         "               return;\n"
 172         "       }\n"
 173         "    int2 sortedData = pHash[index];\n"
 174         "       int unsorted_indx = sortedData.y;\n"
 175         "       float4 bbMin = allpAABB[smallAabbMapping[unsorted_indx]*2 + 0];\n"
 176         "       float4 bbMax = allpAABB[smallAabbMapping[unsorted_indx]*2 + 1];\n"
 177         "       float4 pos;\n"
 178         "       pos.x = (bbMin.x + bbMax.x) * 0.5f;\n"
 179         "       pos.y = (bbMin.y + bbMax.y) * 0.5f;\n"
 180         "       pos.z = (bbMin.z + bbMax.z) * 0.5f;\n"
 181         "    // get address in grid\n"
 182         "    int4 gridPosA = getGridPos(pos, pParams);\n"
 183         "    int4 gridPosB; \n"
 184         "    // examine only neighbouring cells\n"
 185         "    for(int z=-1; z<=1; z++) \n"
 186         "    {\n"
 187         "               gridPosB.z = gridPosA.z + z;\n"
 188         "        for(int y=-1; y<=1; y++) \n"
 189         "        {\n"
 190         "                       gridPosB.y = gridPosA.y + y;\n"
 191         "            for(int x=-1; x<=1; x++) \n"
 192         "            {\n"
 193         "                               gridPosB.x = gridPosA.x + x;\n"
 194         "                findPairsInCell(numObjects, gridPosB, index, pHash, pCellStart, allpAABB,smallAabbMapping, pParams, pairCount,pPairBuff2, maxPairs);\n"
 195         "            }\n"
 196         "        }\n"
 197         "    }\n"
 198         "}\n";