rename gpu_sat -> gpu_narrowphase
This commit is contained in:
erwin coumans
291
opencl/gpu_narrowphase/kernels/bvhTraversal.cl
Normal file
291
opencl/gpu_narrowphase/kernels/bvhTraversal.cl
Normal file
@@ -0,0 +1,291 @@
|
||||
//keep this enum in sync with the CPU version (in btCollidable.h)
|
||||
//written by Erwin Coumans
|
||||
|
||||
#define SHAPE_CONVEX_HULL 3
|
||||
#define SHAPE_CONCAVE_TRIMESH 5
|
||||
#define TRIANGLE_NUM_CONVEX_FACES 5
|
||||
#define SHAPE_COMPOUND_OF_CONVEX_HULLS 6
|
||||
#define SHAPE_SPHERE 7
|
||||
|
||||
typedef unsigned int u32;
|
||||
|
||||
#define MAX_NUM_PARTS_IN_BITS 10
|
||||
|
||||
///btQuantizedBvhNode is a compressed aabb node, 16 bytes.
|
||||
///Node can be used for leafnode or internal node. Leafnodes can point to 32-bit triangle index (non-negative range).
|
||||
typedef struct
|
||||
{
|
||||
//12 bytes
|
||||
unsigned short int m_quantizedAabbMin[3];
|
||||
unsigned short int m_quantizedAabbMax[3];
|
||||
//4 bytes
|
||||
int m_escapeIndexOrTriangleIndex;
|
||||
} btQuantizedBvhNode;
|
||||
/*
|
||||
bool isLeafNode() const
|
||||
{
|
||||
//skipindex is negative (internal node), triangleindex >=0 (leafnode)
|
||||
return (m_escapeIndexOrTriangleIndex >= 0);
|
||||
}
|
||||
int getEscapeIndex() const
|
||||
{
|
||||
btAssert(!isLeafNode());
|
||||
return -m_escapeIndexOrTriangleIndex;
|
||||
}
|
||||
int getTriangleIndex() const
|
||||
{
|
||||
btAssert(isLeafNode());
|
||||
unsigned int x=0;
|
||||
unsigned int y = (~(x&0))<<(31-MAX_NUM_PARTS_IN_BITS);
|
||||
// Get only the lower bits where the triangle index is stored
|
||||
return (m_escapeIndexOrTriangleIndex&~(y));
|
||||
}
|
||||
int getPartId() const
|
||||
{
|
||||
btAssert(isLeafNode());
|
||||
// Get only the highest bits where the part index is stored
|
||||
return (m_escapeIndexOrTriangleIndex>>(31-MAX_NUM_PARTS_IN_BITS));
|
||||
}
|
||||
*/
|
||||
|
||||
int getTriangleIndex(const btQuantizedBvhNode* rootNode)
|
||||
{
|
||||
unsigned int x=0;
|
||||
unsigned int y = (~(x&0))<<(31-MAX_NUM_PARTS_IN_BITS);
|
||||
// Get only the lower bits where the triangle index is stored
|
||||
return (rootNode->m_escapeIndexOrTriangleIndex&~(y));
|
||||
}
|
||||
|
||||
int isLeaf(const btQuantizedBvhNode* rootNode)
|
||||
{
|
||||
//skipindex is negative (internal node), triangleindex >=0 (leafnode)
|
||||
return (rootNode->m_escapeIndexOrTriangleIndex >= 0)? 1 : 0;
|
||||
}
|
||||
|
||||
int getEscapeIndex(const btQuantizedBvhNode* rootNode)
|
||||
{
|
||||
return -rootNode->m_escapeIndexOrTriangleIndex;
|
||||
}
|
||||
|
||||
typedef struct
|
||||
{
|
||||
//12 bytes
|
||||
unsigned short int m_quantizedAabbMin[3];
|
||||
unsigned short int m_quantizedAabbMax[3];
|
||||
//4 bytes, points to the root of the subtree
|
||||
int m_rootNodeIndex;
|
||||
//4 bytes
|
||||
int m_subtreeSize;
|
||||
int m_padding[3];
|
||||
} btBvhSubtreeInfo;
|
||||
|
||||
///keep this in sync with btCollidable.h
|
||||
typedef struct
|
||||
{
|
||||
int m_numChildShapes;
|
||||
int blaat2;
|
||||
int m_shapeType;
|
||||
int m_shapeIndex;
|
||||
|
||||
} btCollidableGpu;
|
||||
|
||||
typedef struct
|
||||
{
|
||||
float4 m_childPosition;
|
||||
float4 m_childOrientation;
|
||||
int m_shapeIndex;
|
||||
int m_unused0;
|
||||
int m_unused1;
|
||||
int m_unused2;
|
||||
} btGpuChildShape;
|
||||
|
||||
|
||||
typedef struct
|
||||
{
|
||||
float4 m_pos;
|
||||
float4 m_quat;
|
||||
float4 m_linVel;
|
||||
float4 m_angVel;
|
||||
|
||||
u32 m_collidableIdx;
|
||||
float m_invMass;
|
||||
float m_restituitionCoeff;
|
||||
float m_frictionCoeff;
|
||||
} BodyData;
|
||||
|
||||
typedef struct
|
||||
{
|
||||
union
|
||||
{
|
||||
float4 m_min;
|
||||
float m_minElems[4];
|
||||
int m_minIndices[4];
|
||||
};
|
||||
union
|
||||
{
|
||||
float4 m_max;
|
||||
float m_maxElems[4];
|
||||
int m_maxIndices[4];
|
||||
};
|
||||
} btAabbCL;
|
||||
|
||||
|
||||
int testQuantizedAabbAgainstQuantizedAabb(
|
||||
const unsigned short int* aabbMin1,
|
||||
const unsigned short int* aabbMax1,
|
||||
const unsigned short int* aabbMin2,
|
||||
const unsigned short int* aabbMax2)
|
||||
{
|
||||
//int overlap = 1;
|
||||
if (aabbMin1[0] > aabbMax2[0])
|
||||
return 0;
|
||||
if (aabbMax1[0] < aabbMin2[0])
|
||||
return 0;
|
||||
if (aabbMin1[1] > aabbMax2[1])
|
||||
return 0;
|
||||
if (aabbMax1[1] < aabbMin2[1])
|
||||
return 0;
|
||||
if (aabbMin1[2] > aabbMax2[2])
|
||||
return 0;
|
||||
if (aabbMax1[2] < aabbMin2[2])
|
||||
return 0;
|
||||
return 1;
|
||||
//overlap = ((aabbMin1[0] > aabbMax2[0]) || (aabbMax1[0] < aabbMin2[0])) ? 0 : overlap;
|
||||
//overlap = ((aabbMin1[2] > aabbMax2[2]) || (aabbMax1[2] < aabbMin2[2])) ? 0 : overlap;
|
||||
//overlap = ((aabbMin1[1] > aabbMax2[1]) || (aabbMax1[1] < aabbMin2[1])) ? 0 : overlap;
|
||||
//return overlap;
|
||||
}
|
||||
|
||||
|
||||
void quantizeWithClamp(unsigned short* out, float4 point2,int isMax, float4 bvhAabbMin, float4 bvhAabbMax, float4 bvhQuantization)
|
||||
{
|
||||
float4 clampedPoint = max(point2,bvhAabbMin);
|
||||
clampedPoint = min (clampedPoint, bvhAabbMax);
|
||||
|
||||
float4 v = (clampedPoint - bvhAabbMin) * bvhQuantization;
|
||||
if (isMax)
|
||||
{
|
||||
out[0] = (unsigned short) (((unsigned short)(v.x+1.f) | 1));
|
||||
out[1] = (unsigned short) (((unsigned short)(v.y+1.f) | 1));
|
||||
out[2] = (unsigned short) (((unsigned short)(v.z+1.f) | 1));
|
||||
} else
|
||||
{
|
||||
out[0] = (unsigned short) (((unsigned short)(v.x) & 0xfffe));
|
||||
out[1] = (unsigned short) (((unsigned short)(v.y) & 0xfffe));
|
||||
out[2] = (unsigned short) (((unsigned short)(v.z) & 0xfffe));
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
|
||||
// work-in-progress
|
||||
__kernel void bvhTraversalKernel( __global const int2* pairs,
|
||||
__global const BodyData* rigidBodies,
|
||||
__global const btCollidableGpu* collidables,
|
||||
__global btAabbCL* aabbs,
|
||||
__global int4* concavePairsOut,
|
||||
__global volatile int* numConcavePairsOut,
|
||||
__global const btBvhSubtreeInfo* subtreeHeaders,
|
||||
__global const btQuantizedBvhNode* quantizedNodes,
|
||||
float4 bvhAabbMin,
|
||||
float4 bvhAabbMax,
|
||||
float4 bvhQuantization,
|
||||
int numSubtreeHeaders,
|
||||
int numPairs,
|
||||
int maxNumConcavePairsCapacity)
|
||||
{
|
||||
int id = get_global_id(0);
|
||||
if (id>=numPairs)
|
||||
return;
|
||||
|
||||
int bodyIndexA = pairs[id].x;
|
||||
int bodyIndexB = pairs[id].y;
|
||||
int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;
|
||||
int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;
|
||||
|
||||
//once the broadphase avoids static-static pairs, we can remove this test
|
||||
if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))
|
||||
{
|
||||
return;
|
||||
}
|
||||
|
||||
if (collidables[collidableIndexA].m_shapeType!=SHAPE_CONCAVE_TRIMESH)
|
||||
return;
|
||||
|
||||
int shapeTypeB = collidables[collidableIndexB].m_shapeType;
|
||||
|
||||
if (shapeTypeB!=SHAPE_CONVEX_HULL &&
|
||||
shapeTypeB!=SHAPE_SPHERE &&
|
||||
shapeTypeB!=SHAPE_COMPOUND_OF_CONVEX_HULLS
|
||||
)
|
||||
return;
|
||||
|
||||
|
||||
unsigned short int quantizedQueryAabbMin[3];
|
||||
unsigned short int quantizedQueryAabbMax[3];
|
||||
quantizeWithClamp(quantizedQueryAabbMin,aabbs[bodyIndexB].m_min,false,bvhAabbMin, bvhAabbMax,bvhQuantization);
|
||||
quantizeWithClamp(quantizedQueryAabbMax,aabbs[bodyIndexB].m_max,true ,bvhAabbMin, bvhAabbMax,bvhQuantization);
|
||||
|
||||
for (int i=0;i<numSubtreeHeaders;i++)
|
||||
{
|
||||
btBvhSubtreeInfo subtree = subtreeHeaders[i];
|
||||
|
||||
int overlap = testQuantizedAabbAgainstQuantizedAabb(quantizedQueryAabbMin,quantizedQueryAabbMax,subtree.m_quantizedAabbMin,subtree.m_quantizedAabbMax);
|
||||
if (overlap != 0)
|
||||
{
|
||||
int startNodeIndex = subtree.m_rootNodeIndex;
|
||||
int endNodeIndex = subtree.m_rootNodeIndex+subtree.m_subtreeSize;
|
||||
int curIndex = startNodeIndex;
|
||||
int escapeIndex;
|
||||
int isLeafNode;
|
||||
int aabbOverlap;
|
||||
while (curIndex < endNodeIndex)
|
||||
{
|
||||
btQuantizedBvhNode rootNode = quantizedNodes[curIndex];
|
||||
aabbOverlap = testQuantizedAabbAgainstQuantizedAabb(quantizedQueryAabbMin,quantizedQueryAabbMax,rootNode.m_quantizedAabbMin,rootNode.m_quantizedAabbMax);
|
||||
isLeafNode = isLeaf(&rootNode);
|
||||
if (aabbOverlap)
|
||||
{
|
||||
if (isLeafNode)
|
||||
{
|
||||
int triangleIndex = getTriangleIndex(&rootNode);
|
||||
if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS)
|
||||
{
|
||||
int numChildrenB = collidables[collidableIndexB].m_numChildShapes;
|
||||
int pairIdx = atomic_add(numConcavePairsOut,numChildrenB);
|
||||
for (int b=0;b<numChildrenB;b++)
|
||||
{
|
||||
if ((pairIdx+b)<maxNumConcavePairsCapacity)
|
||||
{
|
||||
int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+b;
|
||||
int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,childShapeIndexB);
|
||||
concavePairsOut[pairIdx+b] = newPair;
|
||||
}
|
||||
}
|
||||
} else
|
||||
{
|
||||
int pairIdx = atomic_inc(numConcavePairsOut);
|
||||
if (pairIdx<maxNumConcavePairsCapacity)
|
||||
{
|
||||
int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,0);
|
||||
concavePairsOut[pairIdx] = newPair;
|
||||
}
|
||||
}
|
||||
}
|
||||
curIndex++;
|
||||
} else
|
||||
{
|
||||
if (isLeafNode)
|
||||
{
|
||||
curIndex++;
|
||||
} else
|
||||
{
|
||||
escapeIndex = getEscapeIndex(&rootNode);
|
||||
curIndex += escapeIndex;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
}
|
||||
294
opencl/gpu_narrowphase/kernels/bvhTraversal.h
Normal file
294
opencl/gpu_narrowphase/kernels/bvhTraversal.h
Normal file
@@ -0,0 +1,294 @@
|
||||
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
|
||||
static const char* bvhTraversalKernelCL= \
|
||||
"//keep this enum in sync with the CPU version (in btCollidable.h)\n"
|
||||
"//written by Erwin Coumans\n"
|
||||
"\n"
|
||||
"#define SHAPE_CONVEX_HULL 3\n"
|
||||
"#define SHAPE_CONCAVE_TRIMESH 5\n"
|
||||
"#define TRIANGLE_NUM_CONVEX_FACES 5\n"
|
||||
"#define SHAPE_COMPOUND_OF_CONVEX_HULLS 6\n"
|
||||
"#define SHAPE_SPHERE 7\n"
|
||||
"\n"
|
||||
"typedef unsigned int u32;\n"
|
||||
"\n"
|
||||
"#define MAX_NUM_PARTS_IN_BITS 10\n"
|
||||
"\n"
|
||||
"///btQuantizedBvhNode is a compressed aabb node, 16 bytes.\n"
|
||||
"///Node can be used for leafnode or internal node. Leafnodes can point to 32-bit triangle index (non-negative range).\n"
|
||||
"typedef struct\n"
|
||||
"{\n"
|
||||
" //12 bytes\n"
|
||||
" unsigned short int m_quantizedAabbMin[3];\n"
|
||||
" unsigned short int m_quantizedAabbMax[3];\n"
|
||||
" //4 bytes\n"
|
||||
" int m_escapeIndexOrTriangleIndex;\n"
|
||||
"} btQuantizedBvhNode;\n"
|
||||
"/*\n"
|
||||
" bool isLeafNode() const\n"
|
||||
" {\n"
|
||||
" //skipindex is negative (internal node), triangleindex >=0 (leafnode)\n"
|
||||
" return (m_escapeIndexOrTriangleIndex >= 0);\n"
|
||||
" }\n"
|
||||
" int getEscapeIndex() const\n"
|
||||
" {\n"
|
||||
" btAssert(!isLeafNode());\n"
|
||||
" return -m_escapeIndexOrTriangleIndex;\n"
|
||||
" }\n"
|
||||
" int getTriangleIndex() const\n"
|
||||
" {\n"
|
||||
" btAssert(isLeafNode());\n"
|
||||
" unsigned int x=0;\n"
|
||||
" unsigned int y = (~(x&0))<<(31-MAX_NUM_PARTS_IN_BITS);\n"
|
||||
" // Get only the lower bits where the triangle index is stored\n"
|
||||
" return (m_escapeIndexOrTriangleIndex&~(y));\n"
|
||||
" }\n"
|
||||
" int getPartId() const\n"
|
||||
" {\n"
|
||||
" btAssert(isLeafNode());\n"
|
||||
" // Get only the highest bits where the part index is stored\n"
|
||||
" return (m_escapeIndexOrTriangleIndex>>(31-MAX_NUM_PARTS_IN_BITS));\n"
|
||||
" }\n"
|
||||
"*/\n"
|
||||
"\n"
|
||||
"int getTriangleIndex(const btQuantizedBvhNode* rootNode)\n"
|
||||
"{\n"
|
||||
" unsigned int x=0;\n"
|
||||
" unsigned int y = (~(x&0))<<(31-MAX_NUM_PARTS_IN_BITS);\n"
|
||||
" // Get only the lower bits where the triangle index is stored\n"
|
||||
" return (rootNode->m_escapeIndexOrTriangleIndex&~(y));\n"
|
||||
"}\n"
|
||||
"\n"
|
||||
"int isLeaf(const btQuantizedBvhNode* rootNode)\n"
|
||||
"{\n"
|
||||
" //skipindex is negative (internal node), triangleindex >=0 (leafnode)\n"
|
||||
" return (rootNode->m_escapeIndexOrTriangleIndex >= 0)? 1 : 0;\n"
|
||||
"}\n"
|
||||
" \n"
|
||||
"int getEscapeIndex(const btQuantizedBvhNode* rootNode)\n"
|
||||
"{\n"
|
||||
" return -rootNode->m_escapeIndexOrTriangleIndex;\n"
|
||||
"}\n"
|
||||
"\n"
|
||||
"typedef struct\n"
|
||||
"{\n"
|
||||
" //12 bytes\n"
|
||||
" unsigned short int m_quantizedAabbMin[3];\n"
|
||||
" unsigned short int m_quantizedAabbMax[3];\n"
|
||||
" //4 bytes, points to the root of the subtree\n"
|
||||
" int m_rootNodeIndex;\n"
|
||||
" //4 bytes\n"
|
||||
" int m_subtreeSize;\n"
|
||||
" int m_padding[3];\n"
|
||||
"} btBvhSubtreeInfo;\n"
|
||||
"\n"
|
||||
"///keep this in sync with btCollidable.h\n"
|
||||
"typedef struct\n"
|
||||
"{\n"
|
||||
" int m_numChildShapes;\n"
|
||||
" int blaat2;\n"
|
||||
" int m_shapeType;\n"
|
||||
" int m_shapeIndex;\n"
|
||||
" \n"
|
||||
"} btCollidableGpu;\n"
|
||||
"\n"
|
||||
"typedef struct\n"
|
||||
"{\n"
|
||||
" float4 m_childPosition;\n"
|
||||
" float4 m_childOrientation;\n"
|
||||
" int m_shapeIndex;\n"
|
||||
" int m_unused0;\n"
|
||||
" int m_unused1;\n"
|
||||
" int m_unused2;\n"
|
||||
"} btGpuChildShape;\n"
|
||||
"\n"
|
||||
"\n"
|
||||
"typedef struct\n"
|
||||
"{\n"
|
||||
" float4 m_pos;\n"
|
||||
" float4 m_quat;\n"
|
||||
" float4 m_linVel;\n"
|
||||
" float4 m_angVel;\n"
|
||||
"\n"
|
||||
" u32 m_collidableIdx;\n"
|
||||
" float m_invMass;\n"
|
||||
" float m_restituitionCoeff;\n"
|
||||
" float m_frictionCoeff;\n"
|
||||
"} BodyData;\n"
|
||||
"\n"
|
||||
"typedef struct \n"
|
||||
"{\n"
|
||||
" union\n"
|
||||
" {\n"
|
||||
" float4 m_min;\n"
|
||||
" float m_minElems[4];\n"
|
||||
" int m_minIndices[4];\n"
|
||||
" };\n"
|
||||
" union\n"
|
||||
" {\n"
|
||||
" float4 m_max;\n"
|
||||
" float m_maxElems[4];\n"
|
||||
" int m_maxIndices[4];\n"
|
||||
" };\n"
|
||||
"} btAabbCL;\n"
|
||||
"\n"
|
||||
"\n"
|
||||
"int testQuantizedAabbAgainstQuantizedAabb(\n"
|
||||
" const unsigned short int* aabbMin1,\n"
|
||||
" const unsigned short int* aabbMax1,\n"
|
||||
" const unsigned short int* aabbMin2,\n"
|
||||
" const unsigned short int* aabbMax2)\n"
|
||||
"{\n"
|
||||
" //int overlap = 1;\n"
|
||||
" if (aabbMin1[0] > aabbMax2[0])\n"
|
||||
" return 0;\n"
|
||||
" if (aabbMax1[0] < aabbMin2[0])\n"
|
||||
" return 0;\n"
|
||||
" if (aabbMin1[1] > aabbMax2[1])\n"
|
||||
" return 0;\n"
|
||||
" if (aabbMax1[1] < aabbMin2[1])\n"
|
||||
" return 0;\n"
|
||||
" if (aabbMin1[2] > aabbMax2[2])\n"
|
||||
" return 0;\n"
|
||||
" if (aabbMax1[2] < aabbMin2[2])\n"
|
||||
" return 0;\n"
|
||||
" return 1;\n"
|
||||
" //overlap = ((aabbMin1[0] > aabbMax2[0]) || (aabbMax1[0] < aabbMin2[0])) ? 0 : overlap;\n"
|
||||
" //overlap = ((aabbMin1[2] > aabbMax2[2]) || (aabbMax1[2] < aabbMin2[2])) ? 0 : overlap;\n"
|
||||
" //overlap = ((aabbMin1[1] > aabbMax2[1]) || (aabbMax1[1] < aabbMin2[1])) ? 0 : overlap;\n"
|
||||
" //return overlap;\n"
|
||||
"}\n"
|
||||
"\n"
|
||||
"\n"
|
||||
"void quantizeWithClamp(unsigned short* out, float4 point2,int isMax, float4 bvhAabbMin, float4 bvhAabbMax, float4 bvhQuantization)\n"
|
||||
"{\n"
|
||||
" float4 clampedPoint = max(point2,bvhAabbMin);\n"
|
||||
" clampedPoint = min (clampedPoint, bvhAabbMax);\n"
|
||||
"\n"
|
||||
" float4 v = (clampedPoint - bvhAabbMin) * bvhQuantization;\n"
|
||||
" if (isMax)\n"
|
||||
" {\n"
|
||||
" out[0] = (unsigned short) (((unsigned short)(v.x+1.f) | 1));\n"
|
||||
" out[1] = (unsigned short) (((unsigned short)(v.y+1.f) | 1));\n"
|
||||
" out[2] = (unsigned short) (((unsigned short)(v.z+1.f) | 1));\n"
|
||||
" } else\n"
|
||||
" {\n"
|
||||
" out[0] = (unsigned short) (((unsigned short)(v.x) & 0xfffe));\n"
|
||||
" out[1] = (unsigned short) (((unsigned short)(v.y) & 0xfffe));\n"
|
||||
" out[2] = (unsigned short) (((unsigned short)(v.z) & 0xfffe));\n"
|
||||
" }\n"
|
||||
"\n"
|
||||
"}\n"
|
||||
"\n"
|
||||
"\n"
|
||||
"// work-in-progress\n"
|
||||
"__kernel void bvhTraversalKernel( __global const int2* pairs, \n"
|
||||
" __global const BodyData* rigidBodies, \n"
|
||||
" __global const btCollidableGpu* collidables,\n"
|
||||
" __global btAabbCL* aabbs,\n"
|
||||
" __global int4* concavePairsOut,\n"
|
||||
" __global volatile int* numConcavePairsOut,\n"
|
||||
" __global const btBvhSubtreeInfo* subtreeHeaders,\n"
|
||||
" __global const btQuantizedBvhNode* quantizedNodes,\n"
|
||||
" float4 bvhAabbMin,\n"
|
||||
" float4 bvhAabbMax,\n"
|
||||
" float4 bvhQuantization,\n"
|
||||
" int numSubtreeHeaders,\n"
|
||||
" int numPairs,\n"
|
||||
" int maxNumConcavePairsCapacity)\n"
|
||||
"{\n"
|
||||
" int id = get_global_id(0);\n"
|
||||
" if (id>=numPairs)\n"
|
||||
" return;\n"
|
||||
" \n"
|
||||
" int bodyIndexA = pairs[id].x;\n"
|
||||
" int bodyIndexB = pairs[id].y;\n"
|
||||
" int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
|
||||
" int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
|
||||
" \n"
|
||||
" //once the broadphase avoids static-static pairs, we can remove this test\n"
|
||||
" if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))\n"
|
||||
" {\n"
|
||||
" return;\n"
|
||||
" }\n"
|
||||
" \n"
|
||||
" if (collidables[collidableIndexA].m_shapeType!=SHAPE_CONCAVE_TRIMESH)\n"
|
||||
" return;\n"
|
||||
"\n"
|
||||
" int shapeTypeB = collidables[collidableIndexB].m_shapeType;\n"
|
||||
" \n"
|
||||
" if (shapeTypeB!=SHAPE_CONVEX_HULL &&\n"
|
||||
" shapeTypeB!=SHAPE_SPHERE &&\n"
|
||||
" shapeTypeB!=SHAPE_COMPOUND_OF_CONVEX_HULLS\n"
|
||||
" )\n"
|
||||
" return;\n"
|
||||
"\n"
|
||||
" \n"
|
||||
" unsigned short int quantizedQueryAabbMin[3];\n"
|
||||
" unsigned short int quantizedQueryAabbMax[3];\n"
|
||||
" quantizeWithClamp(quantizedQueryAabbMin,aabbs[bodyIndexB].m_min,false,bvhAabbMin, bvhAabbMax,bvhQuantization);\n"
|
||||
" quantizeWithClamp(quantizedQueryAabbMax,aabbs[bodyIndexB].m_max,true ,bvhAabbMin, bvhAabbMax,bvhQuantization);\n"
|
||||
" \n"
|
||||
" for (int i=0;i<numSubtreeHeaders;i++)\n"
|
||||
" {\n"
|
||||
" btBvhSubtreeInfo subtree = subtreeHeaders[i];\n"
|
||||
" \n"
|
||||
" int overlap = testQuantizedAabbAgainstQuantizedAabb(quantizedQueryAabbMin,quantizedQueryAabbMax,subtree.m_quantizedAabbMin,subtree.m_quantizedAabbMax);\n"
|
||||
" if (overlap != 0)\n"
|
||||
" {\n"
|
||||
" int startNodeIndex = subtree.m_rootNodeIndex;\n"
|
||||
" int endNodeIndex = subtree.m_rootNodeIndex+subtree.m_subtreeSize;\n"
|
||||
" int curIndex = startNodeIndex;\n"
|
||||
" int escapeIndex;\n"
|
||||
" int isLeafNode;\n"
|
||||
" int aabbOverlap;\n"
|
||||
" while (curIndex < endNodeIndex)\n"
|
||||
" {\n"
|
||||
" btQuantizedBvhNode rootNode = quantizedNodes[curIndex];\n"
|
||||
" aabbOverlap = testQuantizedAabbAgainstQuantizedAabb(quantizedQueryAabbMin,quantizedQueryAabbMax,rootNode.m_quantizedAabbMin,rootNode.m_quantizedAabbMax);\n"
|
||||
" isLeafNode = isLeaf(&rootNode);\n"
|
||||
" if (aabbOverlap)\n"
|
||||
" {\n"
|
||||
" if (isLeafNode)\n"
|
||||
" {\n"
|
||||
" int triangleIndex = getTriangleIndex(&rootNode);\n"
|
||||
" if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
|
||||
" {\n"
|
||||
" int numChildrenB = collidables[collidableIndexB].m_numChildShapes;\n"
|
||||
" int pairIdx = atomic_add(numConcavePairsOut,numChildrenB);\n"
|
||||
" for (int b=0;b<numChildrenB;b++)\n"
|
||||
" {\n"
|
||||
" if ((pairIdx+b)<maxNumConcavePairsCapacity)\n"
|
||||
" {\n"
|
||||
" int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+b;\n"
|
||||
" int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,childShapeIndexB);\n"
|
||||
" concavePairsOut[pairIdx+b] = newPair;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" } else\n"
|
||||
" {\n"
|
||||
" int pairIdx = atomic_inc(numConcavePairsOut);\n"
|
||||
" if (pairIdx<maxNumConcavePairsCapacity)\n"
|
||||
" {\n"
|
||||
" int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,0);\n"
|
||||
" concavePairsOut[pairIdx] = newPair;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" } \n"
|
||||
" curIndex++;\n"
|
||||
" } else\n"
|
||||
" {\n"
|
||||
" if (isLeafNode)\n"
|
||||
" {\n"
|
||||
" curIndex++;\n"
|
||||
" } else\n"
|
||||
" {\n"
|
||||
" escapeIndex = getEscapeIndex(&rootNode);\n"
|
||||
" curIndex += escapeIndex;\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
" }\n"
|
||||
"\n"
|
||||
"}\n"
|
||||
;
|
||||
1353
opencl/gpu_narrowphase/kernels/primitiveContacts.cl
Normal file
1353
opencl/gpu_narrowphase/kernels/primitiveContacts.cl
Normal file
File diff suppressed because it is too large
Load Diff
1356
opencl/gpu_narrowphase/kernels/primitiveContacts.h
Normal file
1356
opencl/gpu_narrowphase/kernels/primitiveContacts.h
Normal file
File diff suppressed because it is too large
Load Diff
1316
opencl/gpu_narrowphase/kernels/sat.cl
Normal file
1316
opencl/gpu_narrowphase/kernels/sat.cl
Normal file
File diff suppressed because it is too large
Load Diff
1920
opencl/gpu_narrowphase/kernels/satClipHullContacts.cl
Normal file
1920
opencl/gpu_narrowphase/kernels/satClipHullContacts.cl
Normal file
File diff suppressed because it is too large
Load Diff
1924
opencl/gpu_narrowphase/kernels/satClipHullContacts.h
Normal file
1924
opencl/gpu_narrowphase/kernels/satClipHullContacts.h
Normal file
File diff suppressed because it is too large
Load Diff
1320
opencl/gpu_narrowphase/kernels/satKernels.h
Normal file
1320
opencl/gpu_narrowphase/kernels/satKernels.h
Normal file
File diff suppressed because it is too large
Load Diff
Reference in New Issue
Block a user