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

Deterministic simulation for concave and compound collision shapes: added childShapeA/B to b3Contact4 + sort on them

Browse Source 
Use tetrahedra instead of barrel for convex demo (until performance for edge-edge is improved)
Increased #overlapping pair capacity from 12 to 16 / objec
This commit is contained in:
erwincoumans
2013-07-17 22:42:50 -07:00
parent ab125fbb6d
commit 733f9027fb
26 changed files with 358 additions and 70 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp
View File

@@ -64,9 +64,15 @@ void GpuConvexScene::destroyScene()
int GpuConvexScene::createDynamicsObjects(const ConstructionInfo& ci) int GpuConvexScene::createDynamicsObjects(const ConstructionInfo& ci)
{ {
int strideInBytes = 9*sizeof(float); int strideInBytes = 9*sizeof(float);
int numVertices = sizeof(barrel_vertices)/strideInBytes; /*int numVertices = sizeof(barrel_vertices)/strideInBytes;
int numIndices = sizeof(barrel_indices)/sizeof(int); int numIndices = sizeof(barrel_indices)/sizeof(int);
return createDynamicsObjects2(ci,barrel_vertices,numVertices,barrel_indices,numIndices); return createDynamicsObjects2(ci,barrel_vertices,numVertices,barrel_indices,numIndices);
*/
int numVertices = sizeof(tetra_vertices)/strideInBytes;
int numIndices = sizeof(tetra_indices)/sizeof(int);
return createDynamicsObjects2(ci,tetra_vertices,numVertices,tetra_indices,numIndices);
} }
int GpuBoxPlaneScene::createDynamicsObjects(const ConstructionInfo& ci) int GpuBoxPlaneScene::createDynamicsObjects(const ConstructionInfo& ci)

2
Demos3/GpuDemos/rigidbody/GpuRigidBodyDemo.cpp
View File

@@ -110,7 +110,7 @@ void GpuRigidBodyDemo::initPhysics(const ConstructionInfo& ci)
b3Config config; b3Config config;
config.m_maxConvexBodies = b3Max(config.m_maxConvexBodies,ci.arraySizeX*ci.arraySizeY*ci.arraySizeZ+10); config.m_maxConvexBodies = b3Max(config.m_maxConvexBodies,ci.arraySizeX*ci.arraySizeY*ci.arraySizeZ+10);
config.m_maxConvexShapes = config.m_maxConvexBodies; config.m_maxConvexShapes = config.m_maxConvexBodies;
config.m_maxBroadphasePairs = 12*config.m_maxConvexBodies; config.m_maxBroadphasePairs = 16*config.m_maxConvexBodies;
config.m_maxContactCapacity = config.m_maxBroadphasePairs; config.m_maxContactCapacity = config.m_maxBroadphasePairs;

6
src/Bullet3Collision/NarrowPhaseCollision/b3Contact4.h
View File

@@ -34,6 +34,12 @@ B3_ATTRIBUTE_ALIGNED16(struct) b3Contact4
int m_bodyAPtrAndSignBit; int m_bodyAPtrAndSignBit;
int m_bodyBPtrAndSignBit; int m_bodyBPtrAndSignBit;
int m_childIndexA;
int m_childIndexB;
int m_unused1;
int m_unused2;
int getBodyA()const {return abs(m_bodyAPtrAndSignBit);} int getBodyA()const {return abs(m_bodyAPtrAndSignBit);}
int getBodyB()const {return abs(m_bodyBPtrAndSignBit);} int getBodyB()const {return abs(m_bodyBPtrAndSignBit);}
bool isBodyAFixed()const { return m_bodyAPtrAndSignBit<0;} bool isBodyAFixed()const { return m_bodyAPtrAndSignBit<0;}

59
src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.cl
View File

@@ -49,6 +49,12 @@ typedef struct
int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr
int m_bodyBPtrAndSignBit; int m_bodyBPtrAndSignBit;
int m_childIndexA;
int m_childIndexB;
int m_unused1;
int m_unused2;
} Contact4; } Contact4;
typedef struct typedef struct
@@ -483,6 +489,9 @@ void computeContactSphereConvex(int pairIndex,
c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA; c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB; c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
c->m_worldPos[0] = pOnB1; c->m_worldPos[0] = pOnB1;
c->m_childIndexA = -1;
c->m_childIndexB = -1;
GET_NPOINTS(*c) = 1; GET_NPOINTS(*c) = 1;
} }
@@ -705,6 +714,8 @@ void computeContactPlaneConvex(int pairIndex,
c->m_batchIdx = pairIndex; c->m_batchIdx = pairIndex;
c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA; c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB; c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
c->m_childIndexA = -1;
c->m_childIndexB = -1;
switch (numReducedPoints) switch (numReducedPoints)
{ {
@@ -783,6 +794,8 @@ void computeContactPlaneSphere(int pairIndex,
c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA; c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB; c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
c->m_worldPos[0] = pOnB1; c->m_worldPos[0] = pOnB1;
c->m_childIndexA = -1;
c->m_childIndexB = -1;
GET_NPOINTS(*c) = 1; GET_NPOINTS(*c) = 1;
}//if (dstIdx < numPairs) }//if (dstIdx < numPairs)
}//if (hasCollision) }//if (hasCollision)
@@ -955,6 +968,8 @@ __kernel void primitiveContactsKernel( __global const int2* pairs,
c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA; c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;
c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB; c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;
c->m_worldPos[0] = contactPosB; c->m_worldPos[0] = contactPosB;
c->m_childIndexA = -1;
c->m_childIndexB = -1;
GET_NPOINTS(*c) = 1; GET_NPOINTS(*c) = 1;
}//if (dstIdx < numPairs) }//if (dstIdx < numPairs)
}//if ( len <= (radiusA+radiusB)) }//if ( len <= (radiusA+radiusB))
@@ -1163,7 +1178,8 @@ void computeContactSphereTriangle(int pairIndex,
float4 spherePos2, float4 spherePos2,
float radius, float radius,
float4 pos, float4 pos,
float4 quat float4 quat,
int faceIndex
) )
{ {
@@ -1176,7 +1192,7 @@ void computeContactSphereTriangle(int pairIndex,
float4 closestPnt = (float4)(0, 0, 0, 0); float4 closestPnt = (float4)(0, 0, 0, 0);
float4 hitNormalWorld = (float4)(0, 0, 0, 0); float4 hitNormalWorld = (float4)(0, 0, 0, 0);
float minDist = -1000000.f; float minDist = -1000000.f;
bool bCollide = true; bool bCollide = false;
////////////////////////////////////// //////////////////////////////////////
@@ -1251,11 +1267,12 @@ void computeContactSphereTriangle(int pairIndex,
closestPnt = contactPoint; closestPnt = contactPoint;
float4 contactToCenter = sphereCenter - contactPoint; float4 contactToCenter = sphereCenter - contactPoint;
minDist = length(contactToCenter); minDist = length(contactToCenter);
if (minDist>0.f) if (minDist>FLT_EPSILON)
{ {
hitNormalWorld = normalize(contactToCenter);//*(1./minDist); hitNormalWorld = normalize(contactToCenter);//*(1./minDist);
bCollide = true;
} }
bCollide = true;
} }
@@ -1273,19 +1290,29 @@ void computeContactSphereTriangle(int pairIndex,
{ {
pOnB1.w = actualDepth; pOnB1.w = actualDepth;
int dstIdx; int dstIdx;
AppendInc( nGlobalContactsOut, dstIdx );
if (dstIdx < maxContactCapacity) float lenSqr = dot3F4(normalOnSurfaceB1,normalOnSurfaceB1);
if (lenSqr>FLT_EPSILON)
{ {
__global Contact4* c = &globalContactsOut[dstIdx]; AppendInc( nGlobalContactsOut, dstIdx );
c->m_worldNormal = normalOnSurfaceB1;
c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16); if (dstIdx < maxContactCapacity)
c->m_batchIdx = pairIndex; {
c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA; __global Contact4* c = &globalContactsOut[dstIdx];
c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB; c->m_worldNormal = normalOnSurfaceB1;
c->m_worldPos[0] = pOnB1; c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
GET_NPOINTS(*c) = 1; c->m_batchIdx = pairIndex;
} c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
c->m_worldPos[0] = pOnB1;
c->m_childIndexA = -1;
c->m_childIndexB = faceIndex;
GET_NPOINTS(*c) = 1;
}
}
} }
}//if (hasCollision) }//if (hasCollision)
@@ -1346,7 +1373,7 @@ __kernel void findConcaveSphereContactsKernel( __global int4* concavePairs,
rigidBodies,collidables, rigidBodies,collidables,
verticesA, verticesA,
globalContactsOut, nGlobalContactsOut,maxContactCapacity, globalContactsOut, nGlobalContactsOut,maxContactCapacity,
spherePos,sphereRadius,convexPos,convexOrn); spherePos,sphereRadius,convexPos,convexOrn, f);
return; return;
} }

59
src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.h
View File

@@ -51,6 +51,12 @@ static const char* primitiveContactsKernelsCL= \
"\n" "\n"
" int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr\n" " int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr\n"
" int m_bodyBPtrAndSignBit;\n" " int m_bodyBPtrAndSignBit;\n"
"\n"
" int m_childIndexA;\n"
" int m_childIndexB;\n"
" int m_unused1;\n"
" int m_unused2;\n"
"\n"
"} Contact4;\n" "} Contact4;\n"
"\n" "\n"
"typedef struct \n" "typedef struct \n"
@@ -485,6 +491,9 @@ static const char* primitiveContactsKernelsCL= \
" c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n" " c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n"
" c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n" " c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n"
" c->m_worldPos[0] = pOnB1;\n" " c->m_worldPos[0] = pOnB1;\n"
" c->m_childIndexA = -1;\n"
" c->m_childIndexB = -1;\n"
"\n"
" GET_NPOINTS(*c) = 1;\n" " GET_NPOINTS(*c) = 1;\n"
" } \n" " } \n"
"\n" "\n"
@@ -707,6 +716,8 @@ static const char* primitiveContactsKernelsCL= \
" c->m_batchIdx = pairIndex;\n" " c->m_batchIdx = pairIndex;\n"
" c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n" " c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n"
" c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n" " c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n"
" c->m_childIndexA = -1;\n"
" c->m_childIndexB = -1;\n"
"\n" "\n"
" switch (numReducedPoints)\n" " switch (numReducedPoints)\n"
" {\n" " {\n"
@@ -785,6 +796,8 @@ static const char* primitiveContactsKernelsCL= \
" c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n" " c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n"
" c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n" " c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n"
" c->m_worldPos[0] = pOnB1;\n" " c->m_worldPos[0] = pOnB1;\n"
" c->m_childIndexA = -1;\n"
" c->m_childIndexB = -1;\n"
" GET_NPOINTS(*c) = 1;\n" " GET_NPOINTS(*c) = 1;\n"
" }//if (dstIdx < numPairs)\n" " }//if (dstIdx < numPairs)\n"
" }//if (hasCollision)\n" " }//if (hasCollision)\n"
@@ -957,6 +970,8 @@ static const char* primitiveContactsKernelsCL= \
" c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n" " c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n"
" c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n" " c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n"
" c->m_worldPos[0] = contactPosB;\n" " c->m_worldPos[0] = contactPosB;\n"
" c->m_childIndexA = -1;\n"
" c->m_childIndexB = -1;\n"
" GET_NPOINTS(*c) = 1;\n" " GET_NPOINTS(*c) = 1;\n"
" }//if (dstIdx < numPairs)\n" " }//if (dstIdx < numPairs)\n"
" }//if ( len <= (radiusA+radiusB))\n" " }//if ( len <= (radiusA+radiusB))\n"
@@ -1165,7 +1180,8 @@ static const char* primitiveContactsKernelsCL= \
" float4 spherePos2,\n" " float4 spherePos2,\n"
" float radius,\n" " float radius,\n"
" float4 pos,\n" " float4 pos,\n"
" float4 quat\n" " float4 quat,\n"
" int faceIndex\n"
" )\n" " )\n"
"{\n" "{\n"
"\n" "\n"
@@ -1178,7 +1194,7 @@ static const char* primitiveContactsKernelsCL= \
" float4 closestPnt = (float4)(0, 0, 0, 0);\n" " float4 closestPnt = (float4)(0, 0, 0, 0);\n"
" float4 hitNormalWorld = (float4)(0, 0, 0, 0);\n" " float4 hitNormalWorld = (float4)(0, 0, 0, 0);\n"
" float minDist = -1000000.f;\n" " float minDist = -1000000.f;\n"
" bool bCollide = true;\n" " bool bCollide = false;\n"
"\n" "\n"
" \n" " \n"
" //////////////////////////////////////\n" " //////////////////////////////////////\n"
@@ -1253,11 +1269,12 @@ static const char* primitiveContactsKernelsCL= \
" closestPnt = contactPoint;\n" " closestPnt = contactPoint;\n"
" float4 contactToCenter = sphereCenter - contactPoint;\n" " float4 contactToCenter = sphereCenter - contactPoint;\n"
" minDist = length(contactToCenter);\n" " minDist = length(contactToCenter);\n"
" if (minDist>0.f)\n" " if (minDist>FLT_EPSILON)\n"
" {\n" " {\n"
" hitNormalWorld = normalize(contactToCenter);//*(1./minDist);\n" " hitNormalWorld = normalize(contactToCenter);//*(1./minDist);\n"
" bCollide = true;\n"
" }\n" " }\n"
" bCollide = true;\n" " \n"
" }\n" " }\n"
"\n" "\n"
"\n" "\n"
@@ -1275,19 +1292,29 @@ static const char* primitiveContactsKernelsCL= \
" {\n" " {\n"
" pOnB1.w = actualDepth;\n" " pOnB1.w = actualDepth;\n"
" int dstIdx;\n" " int dstIdx;\n"
" AppendInc( nGlobalContactsOut, dstIdx );\n" "\n"
" \n" " \n"
" if (dstIdx < maxContactCapacity)\n" " float lenSqr = dot3F4(normalOnSurfaceB1,normalOnSurfaceB1);\n"
" if (lenSqr>FLT_EPSILON)\n"
" {\n" " {\n"
" __global Contact4* c = &globalContactsOut[dstIdx];\n" " AppendInc( nGlobalContactsOut, dstIdx );\n"
" c->m_worldNormal = normalOnSurfaceB1;\n" " \n"
" c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);\n" " if (dstIdx < maxContactCapacity)\n"
" c->m_batchIdx = pairIndex;\n" " {\n"
" c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n" " __global Contact4* c = &globalContactsOut[dstIdx];\n"
" c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n" " c->m_worldNormal = normalOnSurfaceB1;\n"
" c->m_worldPos[0] = pOnB1;\n" " c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);\n"
" GET_NPOINTS(*c) = 1;\n" " c->m_batchIdx = pairIndex;\n"
" } \n" " c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n"
" c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n"
" c->m_worldPos[0] = pOnB1;\n"
"\n"
" c->m_childIndexA = -1;\n"
" c->m_childIndexB = faceIndex;\n"
"\n"
" GET_NPOINTS(*c) = 1;\n"
" } \n"
" }\n"
"\n" "\n"
" }\n" " }\n"
" }//if (hasCollision)\n" " }//if (hasCollision)\n"
@@ -1348,7 +1375,7 @@ static const char* primitiveContactsKernelsCL= \
" rigidBodies,collidables,\n" " rigidBodies,collidables,\n"
" verticesA,\n" " verticesA,\n"
" globalContactsOut, nGlobalContactsOut,maxContactCapacity,\n" " globalContactsOut, nGlobalContactsOut,maxContactCapacity,\n"
" spherePos,sphereRadius,convexPos,convexOrn);\n" " spherePos,sphereRadius,convexPos,convexOrn, f);\n"
"\n" "\n"
" return;\n" " return;\n"
" }\n" " }\n"

27
src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.cl
View File

@@ -50,6 +50,12 @@ typedef struct
int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr
int m_bodyBPtrAndSignBit; int m_bodyBPtrAndSignBit;
int m_childIndexA;
int m_childIndexB;
int m_unused1;
int m_unused2;
} Contact4; } Contact4;
@@ -924,6 +930,8 @@ __kernel void extractManifoldAndAddContactKernel(__global const int2* pairs,
int bodyB = pairs[pairIndex].y; int bodyB = pairs[pairIndex].y;
c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0 ? -bodyA:bodyA; c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0 ? -bodyA:bodyA;
c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0 ? -bodyB:bodyB; c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0 ? -bodyB:bodyB;
c->m_childIndexA = -1;
c->m_childIndexB = -1;
for (int i=0;i<nContacts;i++) for (int i=0;i<nContacts;i++)
{ {
c->m_worldPos[i] = localPoints[contactIdx[i]]; c->m_worldPos[i] = localPoints[contactIdx[i]];
@@ -1034,6 +1042,8 @@ __kernel void clipHullHullKernel( __global const int2* pairs,
int bodyB = pairs[pairIndex].y; int bodyB = pairs[pairIndex].y;
c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA; c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;
c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB; c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;
c->m_childIndexA = -1;
c->m_childIndexB = -1;
for (int i=0;i<nReducedContacts;i++) for (int i=0;i<nReducedContacts;i++)
{ {
@@ -1165,7 +1175,8 @@ __kernel void clipCompoundsHullHullKernel( __global const int4* gpuCompoundPai
int bodyB = gpuCompoundPairs[pairIndex].y; int bodyB = gpuCompoundPairs[pairIndex].y;
c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA; c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;
c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB; c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;
c->m_childIndexA = childShapeIndexA;
c->m_childIndexB = childShapeIndexB;
for (int i=0;i<nReducedContacts;i++) for (int i=0;i<nReducedContacts;i++)
{ {
c->m_worldPos[i] = pointsIn[contactIdx[i]]; c->m_worldPos[i] = pointsIn[contactIdx[i]];
@@ -1241,6 +1252,9 @@ __kernel void sphereSphereCollisionKernel( __global const int2* pairs,
c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA; c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;
c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB; c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;
c->m_worldPos[0] = contactPosB; c->m_worldPos[0] = contactPosB;
c->m_childIndexA = -1;
c->m_childIndexB = -1;
GET_NPOINTS(*c) = 1; GET_NPOINTS(*c) = 1;
}//if (dstIdx < numPairs) }//if (dstIdx < numPairs)
}//if ( len <= (radiusA+radiusB)) }//if ( len <= (radiusA+radiusB))
@@ -1285,6 +1299,7 @@ __kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
int bodyIndexA = concavePairsIn[i].x; int bodyIndexA = concavePairsIn[i].x;
int bodyIndexB = concavePairsIn[i].y; int bodyIndexB = concavePairsIn[i].y;
int f = concavePairsIn[i].z; int f = concavePairsIn[i].z;
int childShapeIndexA = f;
int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx; int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;
int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx; int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;
@@ -1411,12 +1426,13 @@ __kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
float4 sepAxis = separatingNormals[i]; float4 sepAxis = separatingNormals[i];
int shapeTypeB = collidables[collidableIndexB].m_shapeType; int shapeTypeB = collidables[collidableIndexB].m_shapeType;
int childShapeIndexB =-1;
if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS) if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS)
{ {
/////////////////// ///////////////////
///compound shape support ///compound shape support
int childShapeIndexB = concavePairsIn[pairIndex].w; childShapeIndexB = concavePairsIn[pairIndex].w;
int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex; int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;
shapeIndexB = collidables[childColIndexB].m_shapeIndex; shapeIndexB = collidables[childColIndexB].m_shapeIndex;
float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition; float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;
@@ -1468,7 +1484,8 @@ __kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
int bodyB = concavePairsIn[pairIndex].y; int bodyB = concavePairsIn[pairIndex].y;
c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA; c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;
c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB; c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;
c->m_childIndexA = childShapeIndexA;
c->m_childIndexB = childShapeIndexB;
for (int i=0;i<nReducedContacts;i++) for (int i=0;i<nReducedContacts;i++)
{ {
c->m_worldPos[i] = pointsIn[contactIdx[i]]; c->m_worldPos[i] = pointsIn[contactIdx[i]];
@@ -1888,7 +1905,9 @@ __kernel void newContactReductionKernel( __global const int2* pairs,
int bodyB = pairs[pairIndex].y; int bodyB = pairs[pairIndex].y;
c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA; c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;
c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB; c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;
c->m_childIndexA =-1;
c->m_childIndexB =-1;
switch (nReducedContacts) switch (nReducedContacts)
{ {
case 4: case 4:

27
src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.h
View File

@@ -52,6 +52,12 @@ static const char* satClipKernelsCL= \
"\n" "\n"
" int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr\n" " int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr\n"
" int m_bodyBPtrAndSignBit;\n" " int m_bodyBPtrAndSignBit;\n"
"\n"
" int m_childIndexA;\n"
" int m_childIndexB;\n"
" int m_unused1;\n"
" int m_unused2;\n"
"\n"
"} Contact4;\n" "} Contact4;\n"
"\n" "\n"
"\n" "\n"
@@ -926,6 +932,8 @@ static const char* satClipKernelsCL= \
" int bodyB = pairs[pairIndex].y;\n" " int bodyB = pairs[pairIndex].y;\n"
" c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0 ? -bodyA:bodyA;\n" " c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0 ? -bodyA:bodyA;\n"
" c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0 ? -bodyB:bodyB;\n" " c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0 ? -bodyB:bodyB;\n"
" c->m_childIndexA = -1;\n"
" c->m_childIndexB = -1;\n"
" for (int i=0;i<nContacts;i++)\n" " for (int i=0;i<nContacts;i++)\n"
" {\n" " {\n"
" c->m_worldPos[i] = localPoints[contactIdx[i]];\n" " c->m_worldPos[i] = localPoints[contactIdx[i]];\n"
@@ -1036,6 +1044,8 @@ static const char* satClipKernelsCL= \
" int bodyB = pairs[pairIndex].y;\n" " int bodyB = pairs[pairIndex].y;\n"
" c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n" " c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n"
" c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n" " c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n"
" c->m_childIndexA = -1;\n"
" c->m_childIndexB = -1;\n"
"\n" "\n"
" for (int i=0;i<nReducedContacts;i++)\n" " for (int i=0;i<nReducedContacts;i++)\n"
" {\n" " {\n"
@@ -1167,7 +1177,8 @@ static const char* satClipKernelsCL= \
" int bodyB = gpuCompoundPairs[pairIndex].y;\n" " int bodyB = gpuCompoundPairs[pairIndex].y;\n"
" c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n" " c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n"
" c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n" " c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n"
"\n" " c->m_childIndexA = childShapeIndexA;\n"
" c->m_childIndexB = childShapeIndexB;\n"
" for (int i=0;i<nReducedContacts;i++)\n" " for (int i=0;i<nReducedContacts;i++)\n"
" {\n" " {\n"
" c->m_worldPos[i] = pointsIn[contactIdx[i]];\n" " c->m_worldPos[i] = pointsIn[contactIdx[i]];\n"
@@ -1243,6 +1254,9 @@ static const char* satClipKernelsCL= \
" c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n" " c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n"
" c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n" " c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n"
" c->m_worldPos[0] = contactPosB;\n" " c->m_worldPos[0] = contactPosB;\n"
" c->m_childIndexA = -1;\n"
" c->m_childIndexB = -1;\n"
"\n"
" GET_NPOINTS(*c) = 1;\n" " GET_NPOINTS(*c) = 1;\n"
" }//if (dstIdx < numPairs)\n" " }//if (dstIdx < numPairs)\n"
" }//if ( len <= (radiusA+radiusB))\n" " }//if ( len <= (radiusA+radiusB))\n"
@@ -1287,6 +1301,7 @@ static const char* satClipKernelsCL= \
" int bodyIndexA = concavePairsIn[i].x;\n" " int bodyIndexA = concavePairsIn[i].x;\n"
" int bodyIndexB = concavePairsIn[i].y;\n" " int bodyIndexB = concavePairsIn[i].y;\n"
" int f = concavePairsIn[i].z;\n" " int f = concavePairsIn[i].z;\n"
" int childShapeIndexA = f;\n"
" \n" " \n"
" int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n" " int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
" int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n" " int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
@@ -1413,12 +1428,13 @@ static const char* satClipKernelsCL= \
" float4 sepAxis = separatingNormals[i];\n" " float4 sepAxis = separatingNormals[i];\n"
" \n" " \n"
" int shapeTypeB = collidables[collidableIndexB].m_shapeType;\n" " int shapeTypeB = collidables[collidableIndexB].m_shapeType;\n"
" int childShapeIndexB =-1;\n"
" if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n" " if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
" {\n" " {\n"
" ///////////////////\n" " ///////////////////\n"
" ///compound shape support\n" " ///compound shape support\n"
" \n" " \n"
" int childShapeIndexB = concavePairsIn[pairIndex].w;\n" " childShapeIndexB = concavePairsIn[pairIndex].w;\n"
" int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;\n" " int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;\n"
" shapeIndexB = collidables[childColIndexB].m_shapeIndex;\n" " shapeIndexB = collidables[childColIndexB].m_shapeIndex;\n"
" float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;\n" " float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;\n"
@@ -1470,7 +1486,8 @@ static const char* satClipKernelsCL= \
" int bodyB = concavePairsIn[pairIndex].y;\n" " int bodyB = concavePairsIn[pairIndex].y;\n"
" c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n" " c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n"
" c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n" " c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n"
"\n" " c->m_childIndexA = childShapeIndexA;\n"
" c->m_childIndexB = childShapeIndexB;\n"
" for (int i=0;i<nReducedContacts;i++)\n" " for (int i=0;i<nReducedContacts;i++)\n"
" {\n" " {\n"
" c->m_worldPos[i] = pointsIn[contactIdx[i]];\n" " c->m_worldPos[i] = pointsIn[contactIdx[i]];\n"
@@ -1890,7 +1907,9 @@ static const char* satClipKernelsCL= \
" int bodyB = pairs[pairIndex].y;\n" " int bodyB = pairs[pairIndex].y;\n"
" c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n" " c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n"
" c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n" " c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n"
" \n" " c->m_childIndexA =-1;\n"
" c->m_childIndexB =-1;\n"
"\n"
" switch (nReducedContacts)\n" " switch (nReducedContacts)\n"
" {\n" " {\n"
" case 4:\n" " case 4:\n"

2
src/Bullet3OpenCL/RigidBody/b3Config.h
View File

@@ -29,7 +29,7 @@ struct b3Config
m_maxTriConvexPairCapacity(256*1024) m_maxTriConvexPairCapacity(256*1024)
{ {
m_maxConvexShapes = m_maxConvexBodies; m_maxConvexShapes = m_maxConvexBodies;
m_maxBroadphasePairs = 12*m_maxConvexBodies; m_maxBroadphasePairs = 16*m_maxConvexBodies;
m_maxContactCapacity = m_maxBroadphasePairs; m_maxContactCapacity = m_maxBroadphasePairs;
} }
}; };

60
src/Bullet3OpenCL/RigidBody/b3GpuBatchingPgsSolver.cpp
View File

@@ -3,8 +3,9 @@ bool b3GpuBatchContacts = true;
bool b3GpuSolveConstraint = true; bool b3GpuSolveConstraint = true;
bool gpuRadixSort=true; bool gpuRadixSort=true;
bool gpuSetSortData = true; bool gpuSetSortData = true;
bool gpuSortContacts = true;
bool optionalSortContactsDeterminism = true; bool optionalSortContactsDeterminism = true;
bool gpuSortContactsDeterminism = true;
#include "b3GpuBatchingPgsSolver.h" #include "b3GpuBatchingPgsSolver.h"
#include "Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.h" #include "Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.h"
@@ -61,6 +62,10 @@ struct b3GpuBatchingPgsSolverInternalData
cl_kernel m_setDeterminismSortDataBodyAKernel; cl_kernel m_setDeterminismSortDataBodyAKernel;
cl_kernel m_setDeterminismSortDataBodyBKernel; cl_kernel m_setDeterminismSortDataBodyBKernel;
cl_kernel m_setDeterminismSortDataChildShapeAKernel;
cl_kernel m_setDeterminismSortDataChildShapeBKernel;
class b3RadixSort32CL* m_sort32; class b3RadixSort32CL* m_sort32;
@@ -143,7 +148,9 @@ b3GpuBatchingPgsSolver::b3GpuBatchingPgsSolver(cl_context ctx,cl_device_id devic
cl_program solveFrictionProg= b3OpenCLUtils::compileCLProgramFromString( ctx, device, solveFrictionSource, &pErrNum,additionalMacros, B3_SOLVER_FRICTION_KERNEL_PATH); cl_program solveFrictionProg= b3OpenCLUtils::compileCLProgramFromString( ctx, device, solveFrictionSource, &pErrNum,additionalMacros, B3_SOLVER_FRICTION_KERNEL_PATH);
b3Assert(solveFrictionProg); b3Assert(solveFrictionProg);
cl_program solverSetup2Prog= b3OpenCLUtils::compileCLProgramFromString( ctx, device, solverSetup2Source, &pErrNum,additionalMacros, B3_SOLVER_SETUP2_KERNEL_PATH); //cl_program solverSetup2Prog= b3OpenCLUtils::compileCLProgramFromString( ctx, device, solverSetup2Source, &pErrNum,additionalMacros, B3_SOLVER_SETUP2_KERNEL_PATH);
cl_program solverSetup2Prog= b3OpenCLUtils::compileCLProgramFromString( ctx, device, 0, &pErrNum,additionalMacros, B3_SOLVER_SETUP2_KERNEL_PATH,true);
b3Assert(solverSetup2Prog); b3Assert(solverSetup2Prog);
@@ -168,6 +175,13 @@ b3GpuBatchingPgsSolver::b3GpuBatchingPgsSolver(cl_context ctx,cl_device_id devic
m_data->m_setDeterminismSortDataBodyBKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, solverSetup2Source, "SetDeterminismSortDataBodyB", &pErrNum, solverSetup2Prog,additionalMacros ); m_data->m_setDeterminismSortDataBodyBKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, solverSetup2Source, "SetDeterminismSortDataBodyB", &pErrNum, solverSetup2Prog,additionalMacros );
b3Assert(m_data->m_setDeterminismSortDataBodyBKernel); b3Assert(m_data->m_setDeterminismSortDataBodyBKernel);
m_data->m_setDeterminismSortDataChildShapeAKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, solverSetup2Source, "SetDeterminismSortDataChildShapeA", &pErrNum, solverSetup2Prog,additionalMacros );
b3Assert(m_data->m_setDeterminismSortDataChildShapeAKernel);
m_data->m_setDeterminismSortDataChildShapeBKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, solverSetup2Source, "SetDeterminismSortDataChildShapeB", &pErrNum, solverSetup2Prog,additionalMacros );
b3Assert(m_data->m_setDeterminismSortDataChildShapeBKernel);
m_data->m_reorderContactKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, solverSetup2Source, "ReorderContactKernel", &pErrNum, solverSetup2Prog,additionalMacros ); m_data->m_reorderContactKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, solverSetup2Source, "ReorderContactKernel", &pErrNum, solverSetup2Prog,additionalMacros );
b3Assert(m_data->m_reorderContactKernel); b3Assert(m_data->m_reorderContactKernel);
@@ -234,6 +248,13 @@ b3GpuBatchingPgsSolver::~b3GpuBatchingPgsSolver()
clReleaseKernel( m_data->m_reorderContactKernel); clReleaseKernel( m_data->m_reorderContactKernel);
clReleaseKernel( m_data->m_copyConstraintKernel); clReleaseKernel( m_data->m_copyConstraintKernel);
clReleaseKernel(m_data->m_setDeterminismSortDataBodyAKernel);
clReleaseKernel(m_data->m_setDeterminismSortDataBodyBKernel);
clReleaseKernel(m_data->m_setDeterminismSortDataChildShapeAKernel);
clReleaseKernel(m_data->m_setDeterminismSortDataChildShapeBKernel);
delete m_data; delete m_data;
} }
@@ -444,7 +465,11 @@ static bool sortfnc(const b3SortData& a,const b3SortData& b)
static bool b3ContactCmp(const b3Contact4& p, const b3Contact4& q) static bool b3ContactCmp(const b3Contact4& p, const b3Contact4& q)
{ {
return ((p.m_bodyAPtrAndSignBit<q.m_bodyAPtrAndSignBit) || return ((p.m_bodyAPtrAndSignBit<q.m_bodyAPtrAndSignBit) ||
(p.m_bodyAPtrAndSignBit==q.m_bodyAPtrAndSignBit) && (p.m_bodyBPtrAndSignBit<q.m_bodyBPtrAndSignBit)); ((p.m_bodyAPtrAndSignBit==q.m_bodyAPtrAndSignBit) && (p.m_bodyBPtrAndSignBit<q.m_bodyBPtrAndSignBit)) ||
((p.m_bodyAPtrAndSignBit==q.m_bodyAPtrAndSignBit) && (p.m_bodyBPtrAndSignBit==q.m_bodyBPtrAndSignBit) && p.m_childIndexA<q.m_childIndexA ) ||
((p.m_bodyAPtrAndSignBit==q.m_bodyAPtrAndSignBit) && (p.m_bodyBPtrAndSignBit==q.m_bodyBPtrAndSignBit) && p.m_childIndexA<q.m_childIndexA ) ||
((p.m_bodyAPtrAndSignBit==q.m_bodyAPtrAndSignBit) && (p.m_bodyBPtrAndSignBit==q.m_bodyBPtrAndSignBit) && p.m_childIndexA==q.m_childIndexA && p.m_childIndexB<q.m_childIndexB)
);
} }
@@ -552,7 +577,7 @@ void b3GpuBatchingPgsSolver::solveContacts(int numBodies, cl_mem bodyBuf, cl_mem
if (optionalSortContactsDeterminism) if (optionalSortContactsDeterminism)
{ {
if (gpuSortContacts) if (gpuSortContactsDeterminism)
{ {
B3_PROFILE("GPU Sort contact constraints (determinism)"); B3_PROFILE("GPU Sort contact constraints (determinism)");
@@ -561,9 +586,24 @@ void b3GpuBatchingPgsSolver::solveContacts(int numBodies, cl_mem bodyBuf, cl_mem
m_data->m_pBufContactOutGPU->copyToCL(m_data->m_pBufContactOutGPUCopy->getBufferCL(),numContacts,0,0); m_data->m_pBufContactOutGPU->copyToCL(m_data->m_pBufContactOutGPUCopy->getBufferCL(),numContacts,0,0);
{ {
b3LauncherCL launcher(m_data->m_queue, m_data->m_setDeterminismSortDataBodyAKernel); b3LauncherCL launcher(m_data->m_queue, m_data->m_setDeterminismSortDataChildShapeBKernel);
launcher.setBuffer(m_data->m_pBufContactOutGPUCopy->getBufferCL());
launcher.setBuffer(m_data->m_contactKeyValues->getBufferCL());
launcher.setConst(numContacts);
launcher.launch1D( numContacts, 64 );
}
m_data->m_solverGPU->m_sort32->execute(*m_data->m_contactKeyValues);
{
b3LauncherCL launcher(m_data->m_queue, m_data->m_setDeterminismSortDataChildShapeAKernel);
launcher.setBuffer(m_data->m_pBufContactOutGPUCopy->getBufferCL());
launcher.setBuffer(m_data->m_contactKeyValues->getBufferCL());
launcher.setConst(numContacts);
launcher.launch1D( numContacts, 64 );
}
m_data->m_solverGPU->m_sort32->execute(*m_data->m_contactKeyValues);
{
b3LauncherCL launcher(m_data->m_queue, m_data->m_setDeterminismSortDataBodyBKernel);
launcher.setBuffer(m_data->m_pBufContactOutGPUCopy->getBufferCL()); launcher.setBuffer(m_data->m_pBufContactOutGPUCopy->getBufferCL());
launcher.setBuffer(m_data->m_contactKeyValues->getBufferCL()); launcher.setBuffer(m_data->m_contactKeyValues->getBufferCL());
launcher.setConst(numContacts); launcher.setConst(numContacts);
@@ -573,7 +613,7 @@ void b3GpuBatchingPgsSolver::solveContacts(int numBodies, cl_mem bodyBuf, cl_mem
m_data->m_solverGPU->m_sort32->execute(*m_data->m_contactKeyValues); m_data->m_solverGPU->m_sort32->execute(*m_data->m_contactKeyValues);
{ {
b3LauncherCL launcher(m_data->m_queue, m_data->m_setDeterminismSortDataBodyBKernel); b3LauncherCL launcher(m_data->m_queue, m_data->m_setDeterminismSortDataBodyAKernel);
launcher.setBuffer(m_data->m_pBufContactOutGPUCopy->getBufferCL()); launcher.setBuffer(m_data->m_pBufContactOutGPUCopy->getBufferCL());
launcher.setBuffer(m_data->m_contactKeyValues->getBufferCL()); launcher.setBuffer(m_data->m_contactKeyValues->getBufferCL());
launcher.setConst(numContacts); launcher.setConst(numContacts);
@@ -582,8 +622,6 @@ void b3GpuBatchingPgsSolver::solveContacts(int numBodies, cl_mem bodyBuf, cl_mem
m_data->m_solverGPU->m_sort32->execute(*m_data->m_contactKeyValues); m_data->m_solverGPU->m_sort32->execute(*m_data->m_contactKeyValues);
//__global Contact4* in, __global Contact4* out, __global int2* sortData, int4 cb )
{ {
B3_PROFILE("gpu reorderContactKernel (determinism)"); B3_PROFILE("gpu reorderContactKernel (determinism)");
@@ -886,9 +924,9 @@ void b3GpuBatchingPgsSolver::solveContacts(int numBodies, cl_mem bodyBuf, cl_mem
int simdWidth =numBodies+1;//-1;//64;//-1;//32; int simdWidth =numBodies+1;//-1;//64;//-1;//32;
int numBatches = sortConstraintByBatch( &cpuContacts[0]+offset, n, simdWidth,csCfg.m_staticIdx ,numBodies); // on GPU //int numBatches = sortConstraintByBatch( &cpuContacts[0]+offset, n, simdWidth,csCfg.m_staticIdx ,numBodies); // on GPU
//int numBatches = sortConstraintByBatch2( &cpuContacts[0]+offset, n, simdWidth,csCfg.m_staticIdx ,numBodies); // on GPU //int numBatches = sortConstraintByBatch2( &cpuContacts[0]+offset, n, simdWidth,csCfg.m_staticIdx ,numBodies); // on GPU
//int numBatches = sortConstraintByBatch3( &cpuContacts[0]+offset, n, simdWidth,csCfg.m_staticIdx ,numBodies); // on GPU int numBatches = sortConstraintByBatch3( &cpuContacts[0]+offset, n, simdWidth,csCfg.m_staticIdx ,numBodies); // on GPU

8
src/Bullet3OpenCL/RigidBody/b3Solver.cpp
View File

@@ -513,6 +513,7 @@ struct SolveTask// : public ThreadPool::Task
if (bodyA.m_invMass) if (bodyA.m_invMass)
{ {
b3Assert(usedBodies[aIdx]==0); b3Assert(usedBodies[aIdx]==0);
usedBodies[aIdx]++;
} }
if (m_wgUsedBodies) if (m_wgUsedBodies)
{ {
@@ -537,12 +538,15 @@ struct SolveTask// : public ThreadPool::Task
} }
} }
} }
usedBodies[aIdx]++;
if (bodyB.m_invMass) if (bodyB.m_invMass)
{ {
b3Assert(usedBodies[bIdx]==0); b3Assert(usedBodies[bIdx]==0);
usedBodies[bIdx]++;
} }
usedBodies[bIdx]++;
if( !m_solveFriction ) if( !m_solveFriction )
{ {

6
src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.cl
View File

@@ -73,6 +73,12 @@ typedef struct
int m_bodyA;//sign bit set for fixed objects int m_bodyA;//sign bit set for fixed objects
int m_bodyB; int m_bodyB;
int m_childIndexA;
int m_childIndexB;
int m_unused1;
int m_unused2;
}Contact4; }Contact4;
typedef struct typedef struct

6
src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.h
View File

@@ -75,6 +75,12 @@ static const char* batchingKernelsCL= \
"\n" "\n"
" int m_bodyA;//sign bit set for fixed objects\n" " int m_bodyA;//sign bit set for fixed objects\n"
" int m_bodyB;\n" " int m_bodyB;\n"
"\n"
" int m_childIndexA;\n"
" int m_childIndexB;\n"
" int m_unused1;\n"
" int m_unused2;\n"
"\n"
"}Contact4;\n" "}Contact4;\n"
"\n" "\n"
"typedef struct \n" "typedef struct \n"

6
src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.cl
View File

@@ -74,6 +74,12 @@ typedef struct
int m_bodyAPtrAndSignBit;//sign bit set for fixed objects int m_bodyAPtrAndSignBit;//sign bit set for fixed objects
int m_bodyBPtrAndSignBit; int m_bodyBPtrAndSignBit;
int m_childIndexA;
int m_childIndexB;
int m_unused1;
int m_unused2;
}Contact4; }Contact4;
typedef struct typedef struct

6
src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.h
View File

@@ -76,6 +76,12 @@ static const char* batchingKernelsNewCL= \
"\n" "\n"
" int m_bodyAPtrAndSignBit;//sign bit set for fixed objects\n" " int m_bodyAPtrAndSignBit;//sign bit set for fixed objects\n"
" int m_bodyBPtrAndSignBit;\n" " int m_bodyBPtrAndSignBit;\n"
"\n"
" int m_childIndexA;\n"
" int m_childIndexB;\n"
" int m_unused1;\n"
" int m_unused2;\n"
"\n"
"}Contact4;\n" "}Contact4;\n"
"\n" "\n"
"typedef struct \n" "typedef struct \n"

2
src/Bullet3OpenCL/RigidBody/kernels/jointSolver.cl
View File

@@ -357,12 +357,10 @@ __kernel void breakViolatedConstraintsKernel(__global b3GpuGenericConstraint* co
int numRows = numConstraintRows[cid]; int numRows = numConstraintRows[cid];
if (numRows) if (numRows)
{ {
// printf("cid=%d, breakingThreshold =%f\n",cid,breakingThreshold);
for (int i=0;i<numRows;i++) for (int i=0;i<numRows;i++)
{ {
int rowIndex = rowOffsets[cid]+i; int rowIndex = rowOffsets[cid]+i;
float breakingThreshold = constraints[cid].m_breakingImpulseThreshold; float breakingThreshold = constraints[cid].m_breakingImpulseThreshold;
// printf("rows[%d].m_appliedImpulse=%f\n",rowIndex,rows[rowIndex].m_appliedImpulse);
if (fabs(rows[rowIndex].m_appliedImpulse) >= breakingThreshold) if (fabs(rows[rowIndex].m_appliedImpulse) >= breakingThreshold)
{ {
constraints[cid].m_flags =0;//&= ~B3_CONSTRAINT_FLAG_ENABLED; constraints[cid].m_flags =0;//&= ~B3_CONSTRAINT_FLAG_ENABLED;

2
src/Bullet3OpenCL/RigidBody/kernels/jointSolver.h
View File

@@ -359,12 +359,10 @@ static const char* solveConstraintRowsCL= \
" int numRows = numConstraintRows[cid];\n" " int numRows = numConstraintRows[cid];\n"
" if (numRows)\n" " if (numRows)\n"
" {\n" " {\n"
" // printf(\"cid=%d, breakingThreshold =%f\n\",cid,breakingThreshold);\n"
" for (int i=0;i<numRows;i++)\n" " for (int i=0;i<numRows;i++)\n"
" {\n" " {\n"
" int rowIndex = rowOffsets[cid]+i;\n" " int rowIndex = rowOffsets[cid]+i;\n"
" float breakingThreshold = constraints[cid].m_breakingImpulseThreshold;\n" " float breakingThreshold = constraints[cid].m_breakingImpulseThreshold;\n"
" // printf(\"rows[%d].m_appliedImpulse=%f\n\",rowIndex,rows[rowIndex].m_appliedImpulse);\n"
" if (fabs(rows[rowIndex].m_appliedImpulse) >= breakingThreshold)\n" " if (fabs(rows[rowIndex].m_appliedImpulse) >= breakingThreshold)\n"
" {\n" " {\n"
" constraints[cid].m_flags =0;//&= ~B3_CONSTRAINT_FLAG_ENABLED;\n" " constraints[cid].m_flags =0;//&= ~B3_CONSTRAINT_FLAG_ENABLED;\n"

6
src/Bullet3OpenCL/RigidBody/kernels/solveContact.cl
View File

@@ -213,6 +213,12 @@ typedef struct
int m_bodyAPtrAndSignBit; int m_bodyAPtrAndSignBit;
int m_bodyBPtrAndSignBit; int m_bodyBPtrAndSignBit;
int m_childIndexA;
int m_childIndexB;
int m_unused1;
int m_unused2;
} Contact4; } Contact4;
typedef struct typedef struct

6
src/Bullet3OpenCL/RigidBody/kernels/solveContact.h
View File

@@ -215,6 +215,12 @@ static const char* solveContactCL= \
"\n" "\n"
" int m_bodyAPtrAndSignBit;\n" " int m_bodyAPtrAndSignBit;\n"
" int m_bodyBPtrAndSignBit;\n" " int m_bodyBPtrAndSignBit;\n"
" \n"
" int m_childIndexA;\n"
" int m_childIndexB;\n"
" int m_unused1;\n"
" int m_unused2;\n"
"\n"
"} Contact4;\n" "} Contact4;\n"
"\n" "\n"
"typedef struct\n" "typedef struct\n"

6
src/Bullet3OpenCL/RigidBody/kernels/solveFriction.cl
View File

@@ -213,6 +213,12 @@ typedef struct
int m_bodyAPtrAndSignBit; int m_bodyAPtrAndSignBit;
int m_bodyBPtrAndSignBit; int m_bodyBPtrAndSignBit;
int m_childIndexA;
int m_childIndexB;
int m_unused1;
int m_unused2;
} Contact4; } Contact4;
typedef struct typedef struct

6
src/Bullet3OpenCL/RigidBody/kernels/solveFriction.h
View File

@@ -215,6 +215,12 @@ static const char* solveFrictionCL= \
"\n" "\n"
" int m_bodyAPtrAndSignBit;\n" " int m_bodyAPtrAndSignBit;\n"
" int m_bodyBPtrAndSignBit;\n" " int m_bodyBPtrAndSignBit;\n"
"\n"
" int m_childIndexA;\n"
" int m_childIndexB;\n"
" int m_unused1;\n"
" int m_unused2;\n"
"\n"
"} Contact4;\n" "} Contact4;\n"
"\n" "\n"
"typedef struct\n" "typedef struct\n"

6
src/Bullet3OpenCL/RigidBody/kernels/solverSetup.cl
View File

@@ -412,6 +412,12 @@ typedef struct
int m_bodyAPtrAndSignBit; int m_bodyAPtrAndSignBit;
int m_bodyBPtrAndSignBit; int m_bodyBPtrAndSignBit;
int m_childIndexA;
int m_childIndexB;
int m_unused1;
int m_unused2;
} Contact4; } Contact4;
typedef struct typedef struct

6
src/Bullet3OpenCL/RigidBody/kernels/solverSetup.h
View File

@@ -414,6 +414,12 @@ static const char* solverSetupCL= \
"\n" "\n"
" int m_bodyAPtrAndSignBit;\n" " int m_bodyAPtrAndSignBit;\n"
" int m_bodyBPtrAndSignBit;\n" " int m_bodyBPtrAndSignBit;\n"
"\n"
" int m_childIndexA;\n"
" int m_childIndexB;\n"
" int m_unused1;\n"
" int m_unused2;\n"
"\n"
"} Contact4;\n" "} Contact4;\n"
"\n" "\n"
"typedef struct\n" "typedef struct\n"

47
src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.cl
View File

@@ -386,6 +386,12 @@ typedef struct
int m_bodyAPtrAndSignBit; int m_bodyAPtrAndSignBit;
int m_bodyBPtrAndSignBit; int m_bodyBPtrAndSignBit;
int m_childIndexA;
int m_childIndexB;
int m_unused1;
int m_unused2;
} Contact4; } Contact4;
typedef struct typedef struct
@@ -441,22 +447,53 @@ void ReorderContactKernel(__global Contact4* in, __global Contact4* out, __globa
} }
} }
__kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1)))
__kernel void SetDeterminismSortDataChildShapeB(__global Contact4* contactsIn, __global int2* sortDataOut, int nContacts)
__attribute__((reqd_work_group_size(WG_SIZE,1,1)))
void SetDeterminismSortDataBodyA(__global Contact4* contactsIn, __global int2* sortDataOut, int nContacts)
{ {
int gIdx = GET_GLOBAL_IDX; int gIdx = GET_GLOBAL_IDX;
if( gIdx < nContacts ) if( gIdx < nContacts )
{ {
int2 sd; int2 sd;
sd.x = contactsIn[gIdx].m_bodyAPtrAndSignBit; sd.x = contactsIn[gIdx].m_childIndexB;
sd.y = gIdx; sd.y = gIdx;
sortDataOut[gIdx] = sd; sortDataOut[gIdx] = sd;
} }
} }
__kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1)))
void SetDeterminismSortDataChildShapeA(__global Contact4* contactsIn, __global int2* sortDataInOut, int nContacts)
{
int gIdx = GET_GLOBAL_IDX;
if( gIdx < nContacts )
{
int2 sdIn;
sdIn = sortDataInOut[gIdx];
int2 sdOut;
sdOut.x = contactsIn[sdIn.y].m_childIndexA;
sdOut.y = sdIn.y;
sortDataInOut[gIdx] = sdOut;
}
}
__kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1)))
void SetDeterminismSortDataBodyA(__global Contact4* contactsIn, __global int2* sortDataInOut, int nContacts)
{
int gIdx = GET_GLOBAL_IDX;
if( gIdx < nContacts )
{
int2 sdIn;
sdIn = sortDataInOut[gIdx];
int2 sdOut;
sdOut.x = contactsIn[sdIn.y].m_bodyAPtrAndSignBit;
sdOut.y = sdIn.y;
sortDataInOut[gIdx] = sdOut;
}
}
__kernel __kernel
__attribute__((reqd_work_group_size(WG_SIZE,1,1))) __attribute__((reqd_work_group_size(WG_SIZE,1,1)))
void SetDeterminismSortDataBodyB(__global Contact4* contactsIn, __global int2* sortDataInOut, int nContacts) void SetDeterminismSortDataBodyB(__global Contact4* contactsIn, __global int2* sortDataInOut, int nContacts)

47
src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.h
View File

@@ -388,6 +388,12 @@ static const char* solverSetup2CL= \
"\n" "\n"
" int m_bodyAPtrAndSignBit;\n" " int m_bodyAPtrAndSignBit;\n"
" int m_bodyBPtrAndSignBit;\n" " int m_bodyBPtrAndSignBit;\n"
"\n"
" int m_childIndexA;\n"
" int m_childIndexB;\n"
" int m_unused1;\n"
" int m_unused2;\n"
"\n"
"} Contact4;\n" "} Contact4;\n"
"\n" "\n"
"typedef struct\n" "typedef struct\n"
@@ -443,22 +449,53 @@ static const char* solverSetup2CL= \
" }\n" " }\n"
"}\n" "}\n"
"\n" "\n"
"\n" "__kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
"__kernel\n" "void SetDeterminismSortDataChildShapeA(__global Contact4* contactsIn, __global int2* sortDataOut, int nContacts)\n"
"__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
"void SetDeterminismSortDataBodyA(__global Contact4* contactsIn, __global int2* sortDataOut, int nContacts)\n"
"{\n" "{\n"
" int gIdx = GET_GLOBAL_IDX;\n" " int gIdx = GET_GLOBAL_IDX;\n"
"\n" "\n"
" if( gIdx < nContacts )\n" " if( gIdx < nContacts )\n"
" {\n" " {\n"
" int2 sd;\n" " int2 sd;\n"
" sd.x = contactsIn[gIdx].m_bodyAPtrAndSignBit;\n" " sd.x = contactsIn[gIdx].m_childIndexA;\n"
" sd.y = gIdx;\n" " sd.y = gIdx;\n"
" sortDataOut[gIdx] = sd;\n" " sortDataOut[gIdx] = sd;\n"
" }\n" " }\n"
"}\n" "}\n"
"\n" "\n"
"__kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
"void SetDeterminismSortDataChildShapeB(__global Contact4* contactsIn, __global int2* sortDataInOut, int nContacts)\n"
"{\n"
" int gIdx = GET_GLOBAL_IDX;\n"
"\n"
" if( gIdx < nContacts )\n"
" {\n"
" int2 sdIn;\n"
" sdIn = sortDataInOut[gIdx];\n"
" int2 sdOut;\n"
" sdOut.x = contactsIn[sdIn.y].m_childIndexB;\n"
" sdOut.y = sdIn.y;\n"
" sortDataInOut[gIdx] = sdOut;\n"
" }\n"
"}\n"
"\n"
"__kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
"void SetDeterminismSortDataBodyA(__global Contact4* contactsIn, __global int2* sortDataInOut, int nContacts)\n"
"{\n"
" int gIdx = GET_GLOBAL_IDX;\n"
"\n"
" if( gIdx < nContacts )\n"
" {\n"
" int2 sdIn;\n"
" sdIn = sortDataInOut[gIdx];\n"
" int2 sdOut;\n"
" sdOut.x = contactsIn[sdIn.y].m_bodyAPtrAndSignBit;\n"
" sdOut.y = sdIn.y;\n"
" sortDataInOut[gIdx] = sdOut;\n"
" }\n"
"}\n"
"\n"
"\n"
"__kernel\n" "__kernel\n"
"__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n" "__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
"void SetDeterminismSortDataBodyB(__global Contact4* contactsIn, __global int2* sortDataInOut, int nContacts)\n" "void SetDeterminismSortDataBodyB(__global Contact4* contactsIn, __global int2* sortDataInOut, int nContacts)\n"

6
src/Bullet3OpenCL/RigidBody/kernels/solverUtils.cl
View File

@@ -389,6 +389,12 @@ typedef struct
int m_bodyAPtrAndSignBit; int m_bodyAPtrAndSignBit;
int m_bodyBPtrAndSignBit; int m_bodyBPtrAndSignBit;
int m_childIndexA;
int m_childIndexB;
int m_unused1;
int m_unused2;
} Contact4; } Contact4;

6
src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h
View File

@@ -391,6 +391,12 @@ static const char* solverUtilsCL= \
"\n" "\n"
" int m_bodyAPtrAndSignBit;\n" " int m_bodyAPtrAndSignBit;\n"
" int m_bodyBPtrAndSignBit;\n" " int m_bodyBPtrAndSignBit;\n"
"\n"
" int m_childIndexA;\n"
" int m_childIndexB;\n"
" int m_unused1;\n"
" int m_unused2;\n"
"\n"
"} Contact4;\n" "} Contact4;\n"
"\n" "\n"
"\n" "\n"