tweaks to demos, add middle-mouse support,
This commit is contained in:
erwin coumans
@@ -32,6 +32,22 @@ static const char* sapFastCL= \
|
||||
" };\n"
|
||||
"} btAabbCL;\n"
|
||||
"\n"
|
||||
"typedef struct \n"
|
||||
"{\n"
|
||||
" union\n"
|
||||
" {\n"
|
||||
" unsigned int m_key;\n"
|
||||
" unsigned int x;\n"
|
||||
" };\n"
|
||||
"\n"
|
||||
" union\n"
|
||||
" {\n"
|
||||
" unsigned int m_value;\n"
|
||||
" unsigned int y;\n"
|
||||
" \n"
|
||||
" };\n"
|
||||
"}b3SortData;\n"
|
||||
"\n"
|
||||
"\n"
|
||||
"/// conservative test for overlap between two aabbs\n"
|
||||
"bool TestAabbAgainstAabb2(const btAabbCL* aabb1, __local const btAabbCL* aabb2);\n"
|
||||
@@ -48,6 +64,260 @@ static const char* sapFastCL= \
|
||||
" return overlap;\n"
|
||||
"}\n"
|
||||
"\n"
|
||||
"__kernel void computePairsIncremental3dSapKernel( __global const uint2* objectMinMaxIndexGPUaxis0,\n"
|
||||
" __global const uint2* objectMinMaxIndexGPUaxis1,\n"
|
||||
" __global const uint2* objectMinMaxIndexGPUaxis2,\n"
|
||||
" __global const uint2* objectMinMaxIndexGPUaxis0prev,\n"
|
||||
" __global const uint2* objectMinMaxIndexGPUaxis1prev,\n"
|
||||
" __global const uint2* objectMinMaxIndexGPUaxis2prev,\n"
|
||||
" __global const b3SortData* sortedAxisGPU0,\n"
|
||||
" __global const b3SortData* sortedAxisGPU1,\n"
|
||||
" __global const b3SortData* sortedAxisGPU2,\n"
|
||||
" __global const b3SortData* sortedAxisGPU0prev,\n"
|
||||
" __global const b3SortData* sortedAxisGPU1prev,\n"
|
||||
" __global const b3SortData* sortedAxisGPU2prev,\n"
|
||||
" __global int2* addedHostPairsGPU,\n"
|
||||
" __global int2* removedHostPairsGPU,\n"
|
||||
" volatile __global int* addedHostPairsCount,\n"
|
||||
" volatile __global int* removedHostPairsCount,\n"
|
||||
" int maxCapacity,\n"
|
||||
" int numObjects)\n"
|
||||
"{\n"
|
||||
" int i = get_global_id(0);\n"
|
||||
" if (i>=numObjects)\n"
|
||||
" return;\n"
|
||||
"\n"
|
||||
" __global const uint2* objectMinMaxIndexGPU[3][2];\n"
|
||||
" objectMinMaxIndexGPU[0][0]=objectMinMaxIndexGPUaxis0;\n"
|
||||
" objectMinMaxIndexGPU[1][0]=objectMinMaxIndexGPUaxis1;\n"
|
||||
" objectMinMaxIndexGPU[2][0]=objectMinMaxIndexGPUaxis2;\n"
|
||||
" objectMinMaxIndexGPU[0][1]=objectMinMaxIndexGPUaxis0prev;\n"
|
||||
" objectMinMaxIndexGPU[1][1]=objectMinMaxIndexGPUaxis1prev;\n"
|
||||
" objectMinMaxIndexGPU[2][1]=objectMinMaxIndexGPUaxis2prev;\n"
|
||||
"\n"
|
||||
" __global const b3SortData* sortedAxisGPU[3][2];\n"
|
||||
" sortedAxisGPU[0][0] = sortedAxisGPU0;\n"
|
||||
" sortedAxisGPU[1][0] = sortedAxisGPU1;\n"
|
||||
" sortedAxisGPU[2][0] = sortedAxisGPU2;\n"
|
||||
" sortedAxisGPU[0][1] = sortedAxisGPU0prev;\n"
|
||||
" sortedAxisGPU[1][1] = sortedAxisGPU1prev;\n"
|
||||
" sortedAxisGPU[2][1] = sortedAxisGPU2prev;\n"
|
||||
"\n"
|
||||
" int m_currentBuffer = 0;\n"
|
||||
"\n"
|
||||
" for (int axis=0;axis<3;axis++)\n"
|
||||
" {\n"
|
||||
" //int i = checkObjects[a];\n"
|
||||
"\n"
|
||||
" unsigned int curMinIndex = objectMinMaxIndexGPU[axis][m_currentBuffer][i].x;\n"
|
||||
" unsigned int curMaxIndex = objectMinMaxIndexGPU[axis][m_currentBuffer][i].y;\n"
|
||||
" unsigned int prevMinIndex = objectMinMaxIndexGPU[axis][1-m_currentBuffer][i].x;\n"
|
||||
" int dmin = curMinIndex - prevMinIndex;\n"
|
||||
" \n"
|
||||
" unsigned int prevMaxIndex = objectMinMaxIndexGPU[axis][1-m_currentBuffer][i].y;\n"
|
||||
"\n"
|
||||
" int dmax = curMaxIndex - prevMaxIndex;\n"
|
||||
" \n"
|
||||
" for (int otherbuffer = 0;otherbuffer<2;otherbuffer++)\n"
|
||||
" {\n"
|
||||
" if (dmin!=0)\n"
|
||||
" {\n"
|
||||
" int stepMin = dmin<0 ? -1 : 1;\n"
|
||||
" for (int j=prevMinIndex;j!=curMinIndex;j+=stepMin)\n"
|
||||
" {\n"
|
||||
" int otherIndex2 = sortedAxisGPU[axis][otherbuffer][j].y;\n"
|
||||
" int otherIndex = otherIndex2/2;\n"
|
||||
" if (otherIndex!=i)\n"
|
||||
" {\n"
|
||||
" bool otherIsMax = ((otherIndex2&1)!=0);\n"
|
||||
"\n"
|
||||
" if (otherIsMax)\n"
|
||||
" {\n"
|
||||
" \n"
|
||||
" bool overlap = true;\n"
|
||||
"\n"
|
||||
" for (int ax=0;ax<3;ax++)\n"
|
||||
" {\n"
|
||||
" if ((objectMinMaxIndexGPU[ax][m_currentBuffer][i].x > objectMinMaxIndexGPU[ax][m_currentBuffer][otherIndex].y) ||\n"
|
||||
" (objectMinMaxIndexGPU[ax][m_currentBuffer][i].y < objectMinMaxIndexGPU[ax][m_currentBuffer][otherIndex].x))\n"
|
||||
" overlap=false;\n"
|
||||
" }\n"
|
||||
"\n"
|
||||
" // b3Assert(overlap2==overlap);\n"
|
||||
"\n"
|
||||
" bool prevOverlap = true;\n"
|
||||
"\n"
|
||||
" for (int ax=0;ax<3;ax++)\n"
|
||||
" {\n"
|
||||
" if ((objectMinMaxIndexGPU[ax][1-m_currentBuffer][i].x > objectMinMaxIndexGPU[ax][1-m_currentBuffer][otherIndex].y) ||\n"
|
||||
" (objectMinMaxIndexGPU[ax][1-m_currentBuffer][i].y < objectMinMaxIndexGPU[ax][1-m_currentBuffer][otherIndex].x))\n"
|
||||
" prevOverlap=false;\n"
|
||||
" }\n"
|
||||
" \n"
|
||||
"\n"
|
||||
" //b3Assert(overlap==overlap2);\n"
|
||||
" \n"
|
||||
"\n"
|
||||
"\n"
|
||||
" if (dmin<0)\n"
|
||||
" {\n"
|
||||
" if (overlap && !prevOverlap)\n"
|
||||
" {\n"
|
||||
" //add a pair\n"
|
||||
" int2 newPair;\n"
|
||||
" if (i<=otherIndex)\n"
|
||||
" {\n"
|
||||
" newPair.x = i;\n"
|
||||
" newPair.y = otherIndex;\n"
|
||||
" } else\n"
|
||||
" {\n"
|
||||
" newPair.x = otherIndex;\n"
|
||||
" newPair.y = i;\n"
|
||||
" }\n"
|
||||
" \n"
|
||||
" {\n"
|
||||
" int curPair = atomic_inc(addedHostPairsCount);\n"
|
||||
" if (curPair<maxCapacity)\n"
|
||||
" {\n"
|
||||
" \n"
|
||||
" addedHostPairsGPU[curPair].x = newPair.x;\n"
|
||||
" addedHostPairsGPU[curPair].y = newPair.y;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
"\n"
|
||||
" }\n"
|
||||
" } \n"
|
||||
" else\n"
|
||||
" {\n"
|
||||
" if (!overlap && prevOverlap)\n"
|
||||
" {\n"
|
||||
" \n"
|
||||
" //remove a pair\n"
|
||||
" int2 removedPair;\n"
|
||||
" if (i<=otherIndex)\n"
|
||||
" {\n"
|
||||
" removedPair.x = i;\n"
|
||||
" removedPair.y = otherIndex;\n"
|
||||
" } else\n"
|
||||
" {\n"
|
||||
" removedPair.x = otherIndex;\n"
|
||||
" removedPair.y = i;\n"
|
||||
" }\n"
|
||||
" {\n"
|
||||
" int curPair = atomic_inc(removedHostPairsCount);\n"
|
||||
" if (curPair<maxCapacity)\n"
|
||||
" {\n"
|
||||
" \n"
|
||||
" removedHostPairsGPU[curPair].x = removedPair.x;\n"
|
||||
" removedHostPairsGPU[curPair].y = removedPair.y;\n"
|
||||
"\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" }//otherisMax\n"
|
||||
" }//if (dmin<0)\n"
|
||||
" }//if (otherIndex!=i)\n"
|
||||
" }//for (int j=\n"
|
||||
" }\n"
|
||||
" \n"
|
||||
" if (dmax!=0)\n"
|
||||
" {\n"
|
||||
" int stepMax = dmax<0 ? -1 : 1;\n"
|
||||
" for (int j=prevMaxIndex;j!=curMaxIndex;j+=stepMax)\n"
|
||||
" {\n"
|
||||
" int otherIndex2 = sortedAxisGPU[axis][otherbuffer][j].y;\n"
|
||||
" int otherIndex = otherIndex2/2;\n"
|
||||
" if (otherIndex!=i)\n"
|
||||
" {\n"
|
||||
" bool otherIsMin = ((otherIndex2&1)==0);\n"
|
||||
" if (otherIsMin)\n"
|
||||
" {\n"
|
||||
" \n"
|
||||
" bool overlap = true;\n"
|
||||
"\n"
|
||||
" for (int ax=0;ax<3;ax++)\n"
|
||||
" {\n"
|
||||
" if ((objectMinMaxIndexGPU[ax][m_currentBuffer][i].x > objectMinMaxIndexGPU[ax][m_currentBuffer][otherIndex].y) ||\n"
|
||||
" (objectMinMaxIndexGPU[ax][m_currentBuffer][i].y < objectMinMaxIndexGPU[ax][m_currentBuffer][otherIndex].x))\n"
|
||||
" overlap=false;\n"
|
||||
" }\n"
|
||||
" //b3Assert(overlap2==overlap);\n"
|
||||
"\n"
|
||||
" bool prevOverlap = true;\n"
|
||||
"\n"
|
||||
" for (int ax=0;ax<3;ax++)\n"
|
||||
" {\n"
|
||||
" if ((objectMinMaxIndexGPU[ax][1-m_currentBuffer][i].x > objectMinMaxIndexGPU[ax][1-m_currentBuffer][otherIndex].y) ||\n"
|
||||
" (objectMinMaxIndexGPU[ax][1-m_currentBuffer][i].y < objectMinMaxIndexGPU[ax][1-m_currentBuffer][otherIndex].x))\n"
|
||||
" prevOverlap=false;\n"
|
||||
" }\n"
|
||||
" \n"
|
||||
"\n"
|
||||
" if (dmax>0)\n"
|
||||
" {\n"
|
||||
" if (overlap && !prevOverlap)\n"
|
||||
" {\n"
|
||||
" //add a pair\n"
|
||||
" int2 newPair;\n"
|
||||
" if (i<=otherIndex)\n"
|
||||
" {\n"
|
||||
" newPair.x = i;\n"
|
||||
" newPair.y = otherIndex;\n"
|
||||
" } else\n"
|
||||
" {\n"
|
||||
" newPair.x = otherIndex;\n"
|
||||
" newPair.y = i;\n"
|
||||
" }\n"
|
||||
" {\n"
|
||||
" int curPair = atomic_inc(addedHostPairsCount);\n"
|
||||
" if (curPair<maxCapacity)\n"
|
||||
" {\n"
|
||||
" \n"
|
||||
" addedHostPairsGPU[curPair].x = newPair.x;\n"
|
||||
" addedHostPairsGPU[curPair].y = newPair.y;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" \n"
|
||||
" }\n"
|
||||
" } \n"
|
||||
" else\n"
|
||||
" {\n"
|
||||
" if (!overlap && prevOverlap)\n"
|
||||
" {\n"
|
||||
" //if (otherIndex2&1==0) -> min?\n"
|
||||
" //remove a pair\n"
|
||||
" int2 removedPair;\n"
|
||||
" if (i<=otherIndex)\n"
|
||||
" {\n"
|
||||
" removedPair.x = i;\n"
|
||||
" removedPair.y = otherIndex;\n"
|
||||
" } else\n"
|
||||
" {\n"
|
||||
" removedPair.x = otherIndex;\n"
|
||||
" removedPair.y = i;\n"
|
||||
" }\n"
|
||||
" {\n"
|
||||
" int curPair = atomic_inc(removedHostPairsCount);\n"
|
||||
" if (curPair<maxCapacity)\n"
|
||||
" {\n"
|
||||
" \n"
|
||||
" removedHostPairsGPU[curPair].x = removedPair.x;\n"
|
||||
" removedHostPairsGPU[curPair].y = removedPair.y;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" \n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" \n"
|
||||
" }//if (dmin<0)\n"
|
||||
" }//if (otherIndex!=i)\n"
|
||||
" }//for (int j=\n"
|
||||
" }\n"
|
||||
" }//for (int otherbuffer\n"
|
||||
" }//for (int axis=0;\n"
|
||||
"\n"
|
||||
"\n"
|
||||
"}\n"
|
||||
"\n"
|
||||
"//computePairsKernelBatchWrite\n"
|
||||
"__kernel void computePairsKernel( __global const btAabbCL* aabbs, volatile __global int2* pairsOut,volatile __global int* pairCount, int numObjects, int axis, int maxPairs)\n"
|
||||
|
||||
Reference in New Issue
Block a user