Extras/RigidBodyGpuPipeline/opencl/broadphase_benchmark/computeAabbKernelOCL.cl

   1 MSTRINGIFY(
   2
   3 typedef struct
   4 {
   5         int bla;
   6         int numElem;
   7 } MyAabbConstDataCL ;
   8
   9 typedef struct
  10 {
  11         float                   minfx;
  12         float                   minfy;
  13         float                   minfz;
  14         unsigned int    index0;
  15         float                   maxfx;
  16         float                   maxfy;
  17         float                   maxfz;
  18         unsigned int    index1;
  19 } btAabbCL;
  20
  21
  22 __kernel void   computeAabb( __global btAabbCL* aabbs,__global float4* positions, MyAabbConstDataCL cb)
  23 {
  24         int nodeID = get_global_id(0);
  25
  26         if( nodeID < cb.numElem )
  27         {
  28                 aabbs[nodeID].minfx = positions[nodeID].x -1.f;
  29                 aabbs[nodeID].minfy = positions[nodeID].y -1.f;
  30                 aabbs[nodeID].minfz = positions[nodeID].z -1.f;
  31                 aabbs[nodeID].index0 = nodeID;
  32                 aabbs[nodeID].maxfx = positions[nodeID].x +1.f;
  33                 aabbs[nodeID].maxfy = positions[nodeID].y +1.f;
  34                 aabbs[nodeID].maxfz = positions[nodeID].z +1.f;
  35                 aabbs[nodeID].index1 = nodeID;
  36         }
  37 }
  38
  39
  40 __kernel void countOverlappingpairs(    int numObjects,
  41                                                                                 __global int* pPairBuff,
  42                                                                                 __global int2* pPairBuffStartCurr,
  43                                                                                 __global int* pPairScan,
  44                                                                                 __global float4* pAABB )
  45 {
  46     int index = get_global_id(0);
  47     if(index >= numObjects)
  48         {
  49                 return;
  50         }
  51         float4 bbMin = pAABB[index * 2];
  52         int handleIndex = as_int(bbMin.w);
  53         int2 start_curr = pPairBuffStartCurr[handleIndex];
  54         int start = start_curr.x;
  55         int curr = start_curr.y;
  56         __global int *pInp = pPairBuff + start;
  57         int num_changes = 0;
  58         for(int k = 0; k < curr; k++, pInp++)
  59         {
  60                 if(((*pInp) & 0x60000000))//either new or existing pairs (ignore old non-overlapping pairs)
  61                 {
  62                         num_changes++;
  63                 }
  64         }
  65         pPairScan[index+1] = num_changes;
  66 }
  67
  68
  69 __kernel void squeezePairCaches(        int numObjects,
  70                                                                                         __global int* pPairBuff,
  71                                                                                         __global int2* pPairBuffStartCurr,
  72                                                                                         __global int* pPairScan,
  73                                                                                         __global int2* pPairOut,
  74                                                                                         __global float4* pAABB )
  75 {
  76     int index = get_global_id(0);
  77     if(index >= numObjects)
  78         {
  79                 return;
  80         }
  81         float4 bbMin = pAABB[index * 2];
  82         int handleIndex = as_int(bbMin.w);
  83         int2 start_curr = pPairBuffStartCurr[handleIndex];
  84         int start = start_curr.x;
  85         int curr = start_curr.y;
  86         __global int* pInp = pPairBuff + start;
  87         __global int2* pOut = pPairOut + pPairScan[index+1];
  88         __global int* pOut2 = pInp;
  89         int num = 0;
  90         for(int k = 0; k < curr; k++, pInp++)
  91         {
  92                 if(((*pInp) & 0x60000000))
  93                 {
  94                         int2    newpair;
  95                         newpair.x = handleIndex;
  96                         newpair.y = (*pInp) & (~0x60000000);
  97                         *pOut = newpair;
  98                         pOut++;
  99                 }
 100                 if((*pInp) & 0x60000000)
 101                 {
 102                         *pOut2 = (*pInp) & (~0x60000000);
 103                         pOut2++;
 104                         num++;
 105                 }
 106         }
 107         int2 newStartCurr;
 108         newStartCurr.x = start;
 109         newStartCurr.y = num;
 110         pPairBuffStartCurr[handleIndex] = newStartCurr;
 111 }
 112 );