Bart/bullet3
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

add support for compound hulls against concave triangle mesh

Browse Source
This commit is contained in:
erwin coumans
2013-04-05 18:29:58 -07:00
parent 358f4f97a2
commit c8fcf779bb
12 changed files with 361 additions and 82 deletions
Download Patch File Download Diff File Expand all files Collapse all files

29
opencl/gpu_sat/kernels/bvhTraversal.cl
View File

@@ -215,7 +215,9 @@ __kernel void bvhTraversalKernel( __global const int2* pairs,
int shapeTypeB = collidables[collidableIndexB].m_shapeType;
if (shapeTypeB!=SHAPE_CONVEX_HULL &&
shapeTypeB!=SHAPE_SPHERE )
shapeTypeB!=SHAPE_SPHERE &&
shapeTypeB!=SHAPE_COMPOUND_OF_CONVEX_HULLS
)
return;
@@ -247,12 +249,27 @@ __kernel void bvhTraversalKernel( __global const int2* pairs,
if (isLeafNode)
{
int triangleIndex = getTriangleIndex(&rootNode);
int pairIdx = atomic_inc(numConcavePairsOut);
if (pairIdx<maxNumConcavePairsCapacity)
if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS)
{
int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,3);
concavePairsOut[pairIdx] = newPair;
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++;

29
opencl/gpu_sat/kernels/bvhTraversal.h
View File

@@ -217,7 +217,9 @@ static const char* bvhTraversalKernelCL= \
" int shapeTypeB = collidables[collidableIndexB].m_shapeType;\n"
" \n"
" if (shapeTypeB!=SHAPE_CONVEX_HULL &&\n"
" shapeTypeB!=SHAPE_SPHERE )\n"
" shapeTypeB!=SHAPE_SPHERE &&\n"
" shapeTypeB!=SHAPE_COMPOUND_OF_CONVEX_HULLS\n"
" )\n"
" return;\n"
"\n"
" \n"
@@ -249,12 +251,27 @@ static const char* bvhTraversalKernelCL= \
" if (isLeafNode)\n"
" {\n"
" int triangleIndex = getTriangleIndex(&rootNode);\n"
" \n"
" int pairIdx = atomic_inc(numConcavePairsOut);\n"
" if (pairIdx<maxNumConcavePairsCapacity)\n"
" if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
" {\n"
" int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,3);\n"
" concavePairsOut[pairIdx] = newPair;\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"

75
opencl/gpu_sat/kernels/sat.cl
View File

@@ -886,7 +886,8 @@ __kernel void findCompoundPairsKernel( __global const int2* pairs,
}//for (int b=0;b<numChildrenB;b++)
return;
}//if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)
if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)
if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONCAVE_TRIMESH)
&& (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS))
{
int numChildrenB = collidables[collidableIndexB].m_numChildShapes;
for (int b=0;b<numChildrenB;b++)
@@ -1059,9 +1060,8 @@ __kernel void findConcaveSeparatingAxisKernel( __global int4* concavePairs,
__global const float4* uniqueEdges,
__global const btGpuFace* faces,
__global const int* indices,
__global const btGpuChildShape* gpuChildShapes,
__global btAabbCL* aabbs,
__global volatile float4* separatingNormals,
__global volatile int* hasSeparatingAxis,
__global float4* concaveSeparatingNormalsOut,
int numConcavePairs
)
@@ -1081,9 +1081,10 @@ __kernel void findConcaveSeparatingAxisKernel( __global int4* concavePairs,
int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;
int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;
if (collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL)
if (collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL&&
collidables[collidableIndexB].m_shapeType!=SHAPE_COMPOUND_OF_CONVEX_HULLS)
{
concavePairs[pairIdx].w = 0;
concavePairs[pairIdx].w = -1;
return;
}
@@ -1229,36 +1230,60 @@ __kernel void findConcaveSeparatingAxisKernel( __global int4* concavePairs,
float4 c1 = transform(&c1local,&posB,&ornB);
const float4 DeltaC2 = c0 - c1;
bool sepA = findSeparatingAxisLocalA( &convexPolyhedronA, &convexShapes[shapeIndexB],rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,
rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,
DeltaC2,
verticesA,uniqueEdgesA,facesA,indicesA,
vertices,uniqueEdges,faces,indices,
&sepAxis,&dmin);
///////////////////
///compound shape support
if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)
{
int compoundChild = concavePairs[pairIdx].w;
int childShapeIndexB = compoundChild;//collidables[collidableIndexB].m_shapeIndex+compoundChild;
int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;
float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;
float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;
float4 newPosB = transform(&childPosB,&posB,&ornB);
float4 newOrnB = qtMul(ornB,childOrnB);
posB = newPosB;
ornB = newOrnB;
shapeIndexB = collidables[childColIndexB].m_shapeIndex;
}
//////////////////
bool sepA = findSeparatingAxisLocalA( &convexPolyhedronA, &convexShapes[shapeIndexB],
posA,ornA,
posB,ornB,
DeltaC2,
verticesA,uniqueEdgesA,facesA,indicesA,
vertices,uniqueEdges,faces,indices,
&sepAxis,&dmin);
hasSeparatingAxis = 4;
if (!sepA)
{
hasSeparatingAxis = 0;
} else
{
bool sepB = findSeparatingAxisLocalB( &convexShapes[shapeIndexB],&convexPolyhedronA,rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,
rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,
DeltaC2,
vertices,uniqueEdges,faces,indices,
verticesA,uniqueEdgesA,facesA,indicesA,
&sepAxis,&dmin);
bool sepB = findSeparatingAxisLocalB( &convexShapes[shapeIndexB],&convexPolyhedronA,
posB,ornB,
posA,ornA,
DeltaC2,
vertices,uniqueEdges,faces,indices,
verticesA,uniqueEdgesA,facesA,indicesA,
&sepAxis,&dmin);
if (!sepB)
{
hasSeparatingAxis = 0;
} else
{
bool sepEE = findSeparatingAxisEdgeEdgeLocalA( &convexPolyhedronA, &convexShapes[shapeIndexB],rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,
rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,
DeltaC2,
verticesA,uniqueEdgesA,facesA,indicesA,
vertices,uniqueEdges,faces,indices,
&sepAxis,&dmin);
bool sepEE = findSeparatingAxisEdgeEdgeLocalA( &convexPolyhedronA, &convexShapes[shapeIndexB],
posA,ornA,
posB,ornB,
DeltaC2,
verticesA,uniqueEdgesA,facesA,indicesA,
vertices,uniqueEdges,faces,indices,
&sepAxis,&dmin);
if (!sepEE)
{
@@ -1277,12 +1302,12 @@ __kernel void findConcaveSeparatingAxisKernel( __global int4* concavePairs,
} else
{
//mark this pair as in-active
concavePairs[pairIdx].w = 0;
concavePairs[pairIdx].w = -1;
}
}
else
{
//mark this pair as in-active
concavePairs[pairIdx].w = 0;
concavePairs[pairIdx].w = -1;
}
}

25
opencl/gpu_sat/kernels/satClipHullContacts.cl
View File

@@ -1256,6 +1256,7 @@ __kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
__global const float4* uniqueEdges,
__global const btGpuFace* faces,
__global const int* indices,
__global const btGpuChildShape* gpuChildShapes,
__global const float4* separatingNormals,
__global Contact4* restrict globalContactsOut,
counter32_t nGlobalContactsOut,
@@ -1277,8 +1278,8 @@ __kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
if (i<numConcavePairs)
{
//magic value to detect that pair is invalid
if (concavePairsIn[i].w!=3)
//negative value means that the pair is invalid
if (concavePairsIn[i].w<0)
return;
int bodyIndexA = concavePairsIn[i].x;
@@ -1413,6 +1414,22 @@ __kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
float4 sepAxis = separatingNormals[i];
int shapeTypeB = collidables[collidableIndexB].m_shapeType;
if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS)
{
///////////////////
///compound shape support
int compoundChild = concavePairsIn[pairIndex].w;
int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+compoundChild;
int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;
float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;
float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;
float4 newPosB = transform(&childPosB,&posB,&ornB);
float4 newOrnB = qtMul(ornB,childOrnB);
posB = newPosB;
ornB = newOrnB;
shapeIndexB = collidables[childColIndexB].m_shapeIndex;
}
////////////////////////////////////////
@@ -1420,8 +1437,8 @@ __kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
int numLocalContactsOut = clipHullAgainstHullLocalA(sepAxis,
&convexPolyhedronA, &convexShapes[shapeIndexB],
rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,
rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,
posA,ornA,
posB,ornB,
worldVertsB1,worldVertsB2,capacityWorldVerts,
minDist, maxDist,
&verticesA,&facesA,&indicesA,

25
opencl/gpu_sat/kernels/satClipHullContacts.h
View File

@@ -1258,6 +1258,7 @@ static const char* satClipKernelsCL= \
" __global const float4* uniqueEdges,\n"
" __global const btGpuFace* faces,\n"
" __global const int* indices,\n"
" __global const btGpuChildShape* gpuChildShapes,\n"
" __global const float4* separatingNormals,\n"
" __global Contact4* restrict globalContactsOut,\n"
" counter32_t nGlobalContactsOut,\n"
@@ -1279,8 +1280,8 @@ static const char* satClipKernelsCL= \
"\n"
" if (i<numConcavePairs)\n"
" {\n"
" //magic value to detect that pair is invalid\n"
" if (concavePairsIn[i].w!=3)\n"
" //negative value means that the pair is invalid\n"
" if (concavePairsIn[i].w<0)\n"
" return;\n"
"\n"
" int bodyIndexA = concavePairsIn[i].x;\n"
@@ -1415,6 +1416,22 @@ static const char* satClipKernelsCL= \
"\n"
" float4 sepAxis = separatingNormals[i];\n"
" \n"
" int shapeTypeB = collidables[collidableIndexB].m_shapeType;\n"
" if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
" {\n"
" ///////////////////\n"
" ///compound shape support\n"
" int compoundChild = concavePairsIn[pairIndex].w;\n"
" int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+compoundChild;\n"
" int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;\n"
" float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;\n"
" float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;\n"
" float4 newPosB = transform(&childPosB,&posB,&ornB);\n"
" float4 newOrnB = qtMul(ornB,childOrnB);\n"
" posB = newPosB;\n"
" ornB = newOrnB;\n"
" shapeIndexB = collidables[childColIndexB].m_shapeIndex;\n"
" }\n"
" \n"
" ////////////////////////////////////////\n"
" \n"
@@ -1422,8 +1439,8 @@ static const char* satClipKernelsCL= \
" \n"
" int numLocalContactsOut = clipHullAgainstHullLocalA(sepAxis,\n"
" &convexPolyhedronA, &convexShapes[shapeIndexB],\n"
" rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,\n"
" rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,\n"
" posA,ornA,\n"
" posB,ornB,\n"
" worldVertsB1,worldVertsB2,capacityWorldVerts,\n"
" minDist, maxDist,\n"
" &verticesA,&facesA,&indicesA,\n"

75
opencl/gpu_sat/kernels/satKernels.h
View File

@@ -888,7 +888,8 @@ static const char* satKernelsCL= \
" }//for (int b=0;b<numChildrenB;b++) \n"
" return;\n"
" }//if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
" if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
" if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONCAVE_TRIMESH) \n"
" && (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS))\n"
" {\n"
" int numChildrenB = collidables[collidableIndexB].m_numChildShapes;\n"
" for (int b=0;b<numChildrenB;b++)\n"
@@ -1061,9 +1062,8 @@ static const char* satKernelsCL= \
" __global const float4* uniqueEdges,\n"
" __global const btGpuFace* faces,\n"
" __global const int* indices,\n"
" __global const btGpuChildShape* gpuChildShapes,\n"
" __global btAabbCL* aabbs,\n"
" __global volatile float4* separatingNormals,\n"
" __global volatile int* hasSeparatingAxis,\n"
" __global float4* concaveSeparatingNormalsOut,\n"
" int numConcavePairs\n"
" )\n"
@@ -1083,9 +1083,10 @@ static const char* satKernelsCL= \
" int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;\n"
" int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;\n"
"\n"
" if (collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL)\n"
" if (collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL&&\n"
" collidables[collidableIndexB].m_shapeType!=SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
" {\n"
" concavePairs[pairIdx].w = 0;\n"
" concavePairs[pairIdx].w = -1;\n"
" return;\n"
" }\n"
"\n"
@@ -1231,36 +1232,60 @@ static const char* satKernelsCL= \
" float4 c1 = transform(&c1local,&posB,&ornB);\n"
" const float4 DeltaC2 = c0 - c1;\n"
" \n"
" bool sepA = findSeparatingAxisLocalA( &convexPolyhedronA, &convexShapes[shapeIndexB],rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,\n"
" rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,\n"
" DeltaC2,\n"
" verticesA,uniqueEdgesA,facesA,indicesA,\n"
" vertices,uniqueEdges,faces,indices,\n"
" &sepAxis,&dmin);\n"
"\n"
"\n"
" ///////////////////\n"
" ///compound shape support\n"
"\n"
" if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
" {\n"
" int compoundChild = concavePairs[pairIdx].w;\n"
" int childShapeIndexB = compoundChild;//collidables[collidableIndexB].m_shapeIndex+compoundChild;\n"
" int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;\n"
" float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;\n"
" float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;\n"
" float4 newPosB = transform(&childPosB,&posB,&ornB);\n"
" float4 newOrnB = qtMul(ornB,childOrnB);\n"
" posB = newPosB;\n"
" ornB = newOrnB;\n"
" shapeIndexB = collidables[childColIndexB].m_shapeIndex;\n"
" }\n"
" //////////////////\n"
"\n"
"\n"
" bool sepA = findSeparatingAxisLocalA( &convexPolyhedronA, &convexShapes[shapeIndexB],\n"
" posA,ornA,\n"
" posB,ornB,\n"
" DeltaC2,\n"
" verticesA,uniqueEdgesA,facesA,indicesA,\n"
" vertices,uniqueEdges,faces,indices,\n"
" &sepAxis,&dmin);\n"
" hasSeparatingAxis = 4;\n"
" if (!sepA)\n"
" {\n"
" hasSeparatingAxis = 0;\n"
" } else\n"
" {\n"
" bool sepB = findSeparatingAxisLocalB( &convexShapes[shapeIndexB],&convexPolyhedronA,rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,\n"
" rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,\n"
" DeltaC2,\n"
" vertices,uniqueEdges,faces,indices,\n"
" verticesA,uniqueEdgesA,facesA,indicesA,\n"
" &sepAxis,&dmin);\n"
" bool sepB = findSeparatingAxisLocalB( &convexShapes[shapeIndexB],&convexPolyhedronA,\n"
" posB,ornB,\n"
" posA,ornA,\n"
" DeltaC2,\n"
" vertices,uniqueEdges,faces,indices,\n"
" verticesA,uniqueEdgesA,facesA,indicesA,\n"
" &sepAxis,&dmin);\n"
"\n"
" if (!sepB)\n"
" {\n"
" hasSeparatingAxis = 0;\n"
" } else\n"
" {\n"
" bool sepEE = findSeparatingAxisEdgeEdgeLocalA( &convexPolyhedronA, &convexShapes[shapeIndexB],rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,\n"
" rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,\n"
" DeltaC2,\n"
" verticesA,uniqueEdgesA,facesA,indicesA,\n"
" vertices,uniqueEdges,faces,indices,\n"
" &sepAxis,&dmin);\n"
" bool sepEE = findSeparatingAxisEdgeEdgeLocalA( &convexPolyhedronA, &convexShapes[shapeIndexB],\n"
" posA,ornA,\n"
" posB,ornB,\n"
" DeltaC2,\n"
" verticesA,uniqueEdgesA,facesA,indicesA,\n"
" vertices,uniqueEdges,faces,indices,\n"
" &sepAxis,&dmin);\n"
" \n"
" if (!sepEE)\n"
" {\n"
@@ -1279,13 +1304,13 @@ static const char* satKernelsCL= \
" } else\n"
" { \n"
" //mark this pair as in-active\n"
" concavePairs[pairIdx].w = 0;\n"
" concavePairs[pairIdx].w = -1;\n"
" }\n"
" }\n"
" else\n"
" { \n"
" //mark this pair as in-active\n"
" concavePairs[pairIdx].w = 0;\n"
" concavePairs[pairIdx].w = -1;\n"
" }\n"
"}\n"
"\n"