Extras/RigidBodyGpuPipeline/opencl/3dGridBroadphase/Shared/bt3dGridBroadphaseOCL.cl

   1
   2 MSTRINGIFY(
   3
   4 int getPosHash(int4 gridPos, __global float4* pParams)
   5 {
   6         int4 gridDim = *((__global int4*)(pParams + 1));
   7         gridPos.x &= gridDim.x - 1;
   8         gridPos.y &= gridDim.y - 1;
   9         gridPos.z &= gridDim.z - 1;
  10         int hash = gridPos.z * gridDim.y * gridDim.x + gridPos.y * gridDim.x + gridPos.x;
  11         return hash;
  12 }
  13
  14 int4 getGridPos(float4 worldPos, __global float4* pParams)
  15 {
  16     int4 gridPos;
  17         int4 gridDim = *((__global int4*)(pParams + 1));
  18     gridPos.x = (int)floor(worldPos.x * pParams[0].x) & (gridDim.x - 1);
  19     gridPos.y = (int)floor(worldPos.y * pParams[0].y) & (gridDim.y - 1);
  20     gridPos.z = (int)floor(worldPos.z * pParams[0].z) & (gridDim.z - 1);
  21     return gridPos;
  22 }
  23
  24
  25 // calculate grid hash value for each body using its AABB
  26 __kernel void kCalcHashAABB(int numObjects, __global float4* pAABB, __global int2* pHash, __global float4* pParams GUID_ARG)
  27 {
  28     int index = get_global_id(0);
  29     if(index >= numObjects)
  30         {
  31                 return;
  32         }
  33         float4 bbMin = pAABB[index*2];
  34         float4 bbMax = pAABB[index*2 + 1];
  35         float4 pos;
  36         pos.x = (bbMin.x + bbMax.x) * 0.5f;
  37         pos.y = (bbMin.y + bbMax.y) * 0.5f;
  38         pos.z = (bbMin.z + bbMax.z) * 0.5f;
  39         pos.w = 0.f;
  40     // get address in grid
  41     int4 gridPos = getGridPos(pos, pParams);
  42     int gridHash = getPosHash(gridPos, pParams);
  43     // store grid hash and body index
  44     int2 hashVal;
  45     hashVal.x = gridHash;
  46     hashVal.y = index;
  47     pHash[index] = hashVal;
  48 }
  49
  50 __kernel void kClearCellStart(  int numCells,
  51                                                                 __global int* pCellStart GUID_ARG)
  52 {
  53     int index = get_global_id(0);
  54     if(index >= numCells)
  55         {
  56                 return;
  57         }
  58         pCellStart[index] = -1;
  59 }
  60
  61 __kernel void kFindCellStart(int numObjects, __global int2* pHash, __global int* cellStart GUID_ARG)
  62 {
  63         __local int sharedHash[513];
  64     int index = get_global_id(0);
  65         int2 sortedData;
  66     if(index < numObjects)
  67         {
  68                 sortedData = pHash[index];
  69                 // Load hash data into shared memory so that we can look
  70                 // at neighboring body's hash value without loading
  71                 // two hash values per thread
  72                 sharedHash[get_local_id(0) + 1] = sortedData.x;
  73                 if((index > 0) && (get_local_id(0) == 0))
  74                 {
  75                         // first thread in block must load neighbor body hash
  76                         sharedHash[0] = pHash[index-1].x;
  77                 }
  78         }
  79     barrier(CLK_LOCAL_MEM_FENCE);
  80     if(index < numObjects)
  81         {
  82                 if((index == 0) || (sortedData.x != sharedHash[get_local_id(0)]))
  83                 {
  84                         cellStart[sortedData.x] = index;
  85                 }
  86         }
  87 }
  88
  89 int testAABBOverlap(float4 min0, float4 max0, float4 min1, float4 max1)
  90 {
  91         return  (min0.x <= max1.x)&& (min1.x <= max0.x) &&
  92                         (min0.y <= max1.y)&& (min1.y <= max0.y) &&
  93                         (min0.z <= max1.z)&& (min1.z <= max0.z);
  94 }
  95
  96
  97
  98
  99
 100 void findPairsInCell(   int numObjects,
 101                                                 int4    gridPos,
 102                                                 int    index,
 103                                                 __global int2*  pHash,
 104                                                 __global int*   pCellStart,
 105                                                 __global float4* pAABB,
 106                                                 __global int*   pPairBuff,
 107                                                 __global int2*  pPairBuffStartCurr,
 108                                                 __global float4* pParams)
 109 {
 110         int4 pGridDim = *((__global int4*)(pParams + 1));
 111         int maxBodiesPerCell = pGridDim.w;
 112     int gridHash = getPosHash(gridPos, pParams);
 113     // get start of bucket for this cell
 114     int bucketStart = pCellStart[gridHash];
 115     if (bucketStart == -1)
 116         {
 117         return;   // cell empty
 118         }
 119         // iterate over bodies in this cell
 120     int2 sortedData = pHash[index];
 121         int unsorted_indx = sortedData.y;
 122     float4 min0 = pAABB[unsorted_indx*2 + 0];
 123         float4 max0 = pAABB[unsorted_indx*2 + 1];
 124         int handleIndex =  as_int(min0.w);
 125         int2 start_curr = pPairBuffStartCurr[handleIndex];
 126         int start = start_curr.x;
 127         int curr = start_curr.y;
 128         int2 start_curr_next = pPairBuffStartCurr[handleIndex+1];
 129         int curr_max = start_curr_next.x - start - 1;
 130         int bucketEnd = bucketStart + maxBodiesPerCell;
 131         bucketEnd = (bucketEnd > numObjects) ? numObjects : bucketEnd;
 132         for(int index2 = bucketStart; index2 < bucketEnd; index2++)
 133         {
 134         int2 cellData = pHash[index2];
 135         if (cellData.x != gridHash)
 136         {
 137                         break;   // no longer in same bucket
 138                 }
 139                 int unsorted_indx2 = cellData.y;
 140         if (unsorted_indx2 < unsorted_indx) // check not colliding with self
 141         {
 142                         float4 min1 = pAABB[unsorted_indx2*2 + 0];
 143                         float4 max1 = pAABB[unsorted_indx2*2 + 1];
 144                         if(testAABBOverlap(min0, max0, min1, max1))
 145                         {
 146                                 int handleIndex2 = as_int(min1.w);
 147                                 int k;
 148                                 for(k = 0; k < curr; k++)
 149                                 {
 150                                         int old_pair = pPairBuff[start+k] & (~0x60000000);
 151                                         if(old_pair == handleIndex2)
 152                                         {
 153                                                 pPairBuff[start+k] |= 0x40000000;
 154                                                 break;
 155                                         }
 156                                 }
 157                                 if(k == curr)
 158                                 {
 159                                         if(curr >= curr_max)
 160                                         { // not a good solution, but let's avoid crash
 161                                                 break;
 162                                         }
 163                                         pPairBuff[start+curr] = handleIndex2 | 0x20000000;
 164                                         curr++;
 165                                 }
 166                         }
 167                 }
 168         }
 169         int2 newStartCurr;
 170         newStartCurr.x = start;
 171         newStartCurr.y = curr;
 172         pPairBuffStartCurr[handleIndex] = newStartCurr;
 173     return;
 174 }
 175
 176 __kernel void kFindOverlappingPairs(    int numObjects,
 177                                                                                 __global float4* pAABB,
 178                                                                                 __global int2* pHash,
 179                                                                                 __global int* pCellStart,
 180                                                                                 __global int* pPairBuff,
 181                                                                                 __global int2* pPairBuffStartCurr,
 182                                                                                 __global float4* pParams GUID_ARG)
 183
 184 {
 185     int index = get_global_id(0);
 186     if(index >= numObjects)
 187         {
 188                 return;
 189         }
 190     int2 sortedData = pHash[index];
 191         int unsorted_indx = sortedData.y;
 192         float4 bbMin = pAABB[unsorted_indx*2 + 0];
 193         float4 bbMax = pAABB[unsorted_indx*2 + 1];
 194         float4 pos;
 195         pos.x = (bbMin.x + bbMax.x) * 0.5f;
 196         pos.y = (bbMin.y + bbMax.y) * 0.5f;
 197         pos.z = (bbMin.z + bbMax.z) * 0.5f;
 198     // get address in grid
 199     int4 gridPosA = getGridPos(pos, pParams);
 200     int4 gridPosB;
 201     // examine only neighbouring cells
 202     for(int z=-1; z<=1; z++)
 203     {
 204                 gridPosB.z = gridPosA.z + z;
 205         for(int y=-1; y<=1; y++)
 206         {
 207                         gridPosB.y = gridPosA.y + y;
 208             for(int x=-1; x<=1; x++)
 209             {
 210                                 gridPosB.x = gridPosA.x + x;
 211                 findPairsInCell(numObjects, gridPosB, index, pHash, pCellStart, pAABB, pPairBuff, pPairBuffStartCurr, pParams);
 212             }
 213         }
 214     }
 215 }
 216
 217
 218 __kernel void kFindPairsLarge(  int numObjects,
 219                                                                 __global float4* pAABB,
 220                                                                 __global int2* pHash,
 221                                                                 __global int* pCellStart,
 222                                                                 __global int* pPairBuff,
 223                                                                 __global int2* pPairBuffStartCurr,
 224                                                                 uint numLarge GUID_ARG)
 225 {
 226     int index = get_global_id(0);
 227     if(index >= numObjects)
 228         {
 229                 return;
 230         }
 231     int2 sortedData = pHash[index];
 232         int unsorted_indx = sortedData.y;
 233         float4 min0 = pAABB[unsorted_indx*2 + 0];
 234         float4 max0 = pAABB[unsorted_indx*2 + 1];
 235         int handleIndex =  as_int(min0.w);
 236         int2 start_curr = pPairBuffStartCurr[handleIndex];
 237         int start = start_curr.x;
 238         int curr = start_curr.y;
 239         int2 start_curr_next = pPairBuffStartCurr[handleIndex+1];
 240         int curr_max = start_curr_next.x - start - 1;
 241     for(uint i = 0; i < numLarge; i++)
 242     {
 243                 int indx2 = numObjects + i;
 244                 float4 min1 = pAABB[indx2*2 + 0];
 245                 float4 max1 = pAABB[indx2*2 + 1];
 246                 if(testAABBOverlap(min0, max0, min1, max1))
 247                 {
 248                         int k;
 249                         int handleIndex2 =  as_int(min1.w);
 250                         for(k = 0; k < curr; k++)
 251                         {
 252                                 int old_pair = pPairBuff[start+k] & (~0x60000000);
 253                                 if(old_pair == handleIndex2)
 254                                 {
 255                                         pPairBuff[start+k] |= 0x40000000;
 256                                         break;
 257                                 }
 258                         }
 259                         if(k == curr)
 260                         {
 261                                 pPairBuff[start+curr] = handleIndex2 | 0x20000000;
 262                                 if(curr >= curr_max)
 263                                 { // not a good solution, but let's avoid crash
 264                                         break;
 265                                 }
 266                                 curr++;
 267                         }
 268                 }
 269     }
 270         int2 newStartCurr;
 271         newStartCurr.x = start;
 272         newStartCurr.y = curr;
 273         pPairBuffStartCurr[handleIndex] = newStartCurr;
 274     return;
 275 }
 276
 277 __kernel void kComputePairCacheChanges( int numObjects,
 278                                                                                 __global int* pPairBuff,
 279                                                                                 __global int2* pPairBuffStartCurr,
 280                                                                                 __global int* pPairScan,
 281                                                                                 __global float4* pAABB GUID_ARG)
 282 {
 283     int index = get_global_id(0);
 284     if(index >= numObjects)
 285         {
 286                 return;
 287         }
 288         float4 bbMin = pAABB[index * 2];
 289         int handleIndex = as_int(bbMin.w);
 290         int2 start_curr = pPairBuffStartCurr[handleIndex];
 291         int start = start_curr.x;
 292         int curr = start_curr.y;
 293         __global int *pInp = pPairBuff + start;
 294         int num_changes = 0;
 295         for(int k = 0; k < curr; k++, pInp++)
 296         {
 297                 if(!((*pInp) & 0x40000000))
 298                 {
 299                         num_changes++;
 300                 }
 301         }
 302         pPairScan[index+1] = num_changes;
 303 }
 304
 305 __kernel void kSqueezeOverlappingPairBuff(      int numObjects,
 306                                                                                         __global int* pPairBuff,
 307                                                                                         __global int2* pPairBuffStartCurr,
 308                                                                                         __global int* pPairScan,
 309                                                                                         __global int* pPairOut,
 310                                                                                         __global float4* pAABB GUID_ARG)
 311 {
 312     int index = get_global_id(0);
 313     if(index >= numObjects)
 314         {
 315                 return;
 316         }
 317         float4 bbMin = pAABB[index * 2];
 318         int handleIndex = as_int(bbMin.w);
 319         int2 start_curr = pPairBuffStartCurr[handleIndex];
 320         int start = start_curr.x;
 321         int curr = start_curr.y;
 322         __global int* pInp = pPairBuff + start;
 323         __global int* pOut = pPairOut + pPairScan[index+1];
 324         __global int* pOut2 = pInp;
 325         int num = 0;
 326         for(int k = 0; k < curr; k++, pInp++)
 327         {
 328                 if(!((*pInp) & 0x40000000))
 329                 {
 330                         *pOut = *pInp;
 331                         pOut++;
 332                 }
 333                 if((*pInp) & 0x60000000)
 334                 {
 335                         *pOut2 = (*pInp) & (~0x60000000);
 336                         pOut2++;
 337                         num++;
 338                 }
 339         }
 340         int2 newStartCurr;
 341         newStartCurr.x = start;
 342         newStartCurr.y = num;
 343         pPairBuffStartCurr[handleIndex] = newStartCurr;
 344 }
 345
 346
 347
 348
 349 );