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

bugfix for compound versus concave: compute the DeltaC2 after computing the (child) transforms for proper center

Browse Source
This commit is contained in:
erwin coumans
2013-04-12 13:02:35 -07:00
parent ec0d2ed523
commit 07b89beb39
9 changed files with 75 additions and 56 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
demo/gpudemo/GpuDemo.h
View File

@@ -38,9 +38,9 @@ public:
:useOpenCL(true), :useOpenCL(true),
preferredOpenCLPlatformIndex(-1), preferredOpenCLPlatformIndex(-1),
preferredOpenCLDeviceIndex(-1), preferredOpenCLDeviceIndex(-1),
arraySizeX(1), arraySizeX(10),
arraySizeY(2), arraySizeY(30),
arraySizeZ(1), arraySizeZ(10),
m_useConcaveMesh(false), m_useConcaveMesh(false),
gapX(14.3), gapX(14.3),
gapY(14.0), gapY(14.0),

2
demo/gpudemo/main_opengl3core.cpp
View File

@@ -65,7 +65,7 @@ btAlignedObjectArray<const char*> demoNames;
int selectedDemo = 0; int selectedDemo = 0;
GpuDemo::CreateFunc* allDemos[]= GpuDemo::CreateFunc* allDemos[]=
{ {
ConcaveCompound2Scene::MyCreateFunc, // ConcaveCompound2Scene::MyCreateFunc,
ConcaveCompoundScene::MyCreateFunc, ConcaveCompoundScene::MyCreateFunc,

18
demo/gpudemo/rigidbody/ConcaveScene.cpp
View File

@@ -160,12 +160,12 @@ void ConcaveScene::createConcaveMesh(const ConstructionInfo& ci)
objLoader* objData = new objLoader(); objLoader* objData = new objLoader();
//char* fileName = "data/slopedPlane100.obj"; //char* fileName = "data/slopedPlane100.obj";
//char* fileName = "data/plane100.obj"; //char* fileName = "data/plane100.obj";
char* fileName = "data/plane100.obj"; // char* fileName = "data/plane100.obj";
//char* fileName = "data/teddy.obj";//"plane.obj"; //char* fileName = "data/teddy.obj";//"plane.obj";
// char* fileName = "data/sponza_closed.obj";//"plane.obj"; // char* fileName = "data/sponza_closed.obj";//"plane.obj";
//char* fileName = "data/leoTest1.obj"; //char* fileName = "data/leoTest1.obj";
//char* fileName = "data/samurai_monastry.obj"; char* fileName = "data/samurai_monastry.obj";
// char* fileName = "data/teddy2_VHACD_CHs.obj"; // char* fileName = "data/teddy2_VHACD_CHs.obj";
btVector3 shift(0,0,0);//0,230,80);//150,-100,-120); btVector3 shift(0,0,0);//0,230,80);//150,-100,-120);
@@ -324,7 +324,7 @@ void ConcaveScene::createDynamicObjects(const ConstructionInfo& ci)
float mass = 1; float mass = 1;
//btVector3 position(-2*ci.gapX+i*ci.gapX,25+j*ci.gapY,-2*ci.gapZ+k*ci.gapZ); //btVector3 position(-2*ci.gapX+i*ci.gapX,25+j*ci.gapY,-2*ci.gapZ+k*ci.gapZ);
btVector3 position(-(ci.arraySizeX/2)*CONCAVE_GAPX+i*CONCAVE_GAPX,50+j*CONCAVE_GAPY,-(ci.arraySizeZ/2)*CONCAVE_GAPZ+k*CONCAVE_GAPZ); btVector3 position(-(ci.arraySizeX/2)*CONCAVE_GAPX+i*CONCAVE_GAPX,150+j*CONCAVE_GAPY,-(ci.arraySizeZ/2)*CONCAVE_GAPZ+k*CONCAVE_GAPZ);
btQuaternion orn(0,0,0,1); btQuaternion orn(0,0,0,1);
btVector4 color = colors[curColor]; btVector4 color = colors[curColor];
@@ -359,8 +359,8 @@ void ConcaveCompound2Scene::createDynamicObjects(const ConstructionInfo& ci)
{ {
objLoader* objData = new objLoader(); objLoader* objData = new objLoader();
//char* fileName = "data/teddy2_VHACD_CHs.obj"; char* fileName = "data/teddy2_VHACD_CHs.obj";
char* fileName = "data/cube_offset.obj"; //char* fileName = "data/cube_offset.obj";
btVector3 shift(0,0,0);//0,230,80);//150,-100,-120); btVector3 shift(0,0,0);//0,230,80);//150,-100,-120);
@@ -523,7 +523,7 @@ void ConcaveCompound2Scene::createDynamicObjects(const ConstructionInfo& ci)
childShapes.push_back(child); childShapes.push_back(child);
colIndex = childColIndex; colIndex = childColIndex;
} }
//colIndex= m_data->m_np->registerCompoundShape(&childShapes); colIndex= m_data->m_np->registerCompoundShape(&childShapes);
} }
@@ -549,10 +549,10 @@ void ConcaveCompound2Scene::createDynamicObjects(const ConstructionInfo& ci)
float mass = 1;//j==0? 0.f : 1.f; float mass = 1;//j==0? 0.f : 1.f;
//btVector3 position(i*10*ci.gapX,j*ci.gapY,k*10*ci.gapZ); //btVector3 position(i*10*ci.gapX,j*ci.gapY,k*10*ci.gapZ);
btVector3 position(i*10*ci.gapX,1+j*ci.gapY,k*10*ci.gapZ); btVector3 position(i*10*ci.gapX,50+j*ci.gapY,k*10*ci.gapZ);
// btQuaternion orn(0,0,0,1); // btQuaternion orn(0,0,0,1);
btQuaternion orn(btVector3(0,0,1),1.3); btQuaternion orn(btVector3(0,0,1),1.8);
btVector4 color = colors[curColor]; btVector4 color = colors[curColor];
curColor++; curColor++;
@@ -665,7 +665,7 @@ btVector3 childPositions[3] = {
{ {
float mass = 1;//j==0? 0.f : 1.f; float mass = 1;//j==0? 0.f : 1.f;
btVector3 position(i*ci.gapX,10+j*ci.gapY,k*ci.gapZ); btVector3 position(i*ci.gapX,150+j*ci.gapY,k*ci.gapZ);
//btQuaternion orn(0,0,0,1); //btQuaternion orn(0,0,0,1);
btQuaternion orn(btVector3(1,0,0),0.7); btQuaternion orn(btVector3(1,0,0),0.7);

34
opencl/gpu_sat/host/ConvexHullContact.cpp
View File

@@ -53,12 +53,17 @@ m_totalContactsOut(m_context, m_queue)
m_totalContactsOut.push_back(0); m_totalContactsOut.push_back(0);
cl_int errNum=0; cl_int errNum=0;
bool disableKernelCaching = true;
if (1) if (1)
{ {
const char* src = satKernelsCL; const char* src = satKernelsCL;
cl_program satProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,src,&errNum,"","opencl/gpu_sat/kernels/sat.cl",true);
char flags[1024]={0};
//#ifdef CL_PLATFORM_INTEL
// sprintf(flags,"-g -s \"%s\"","C:/develop/bullet3_experiments2/opencl/gpu_sat/kernels/sat.cl");
//#endif
cl_program satProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,src,&errNum,flags,"opencl/gpu_sat/kernels/sat.cl");
btAssert(errNum==CL_SUCCESS); btAssert(errNum==CL_SUCCESS);
m_findSeparatingAxisKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,src, "findSeparatingAxisKernel",&errNum,satProg ); m_findSeparatingAxisKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,src, "findSeparatingAxisKernel",&errNum,satProg );
@@ -80,7 +85,13 @@ m_totalContactsOut(m_context, m_queue)
if (1) if (1)
{ {
const char* srcClip = satClipKernelsCL; const char* srcClip = satClipKernelsCL;
cl_program satClipContactsProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,srcClip,&errNum,"","opencl/gpu_sat/kernels/satClipHullContacts.cl",true);
char flags[1024]={0};
//#ifdef CL_PLATFORM_INTEL
// sprintf(flags,"-g -s \"%s\"","C:/develop/bullet3_experiments2/opencl/gpu_sat/kernels/satClipHullContacts.cl");
//#endif
cl_program satClipContactsProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,srcClip,&errNum,flags,"opencl/gpu_sat/kernels/satClipHullContacts.cl");
btAssert(errNum==CL_SUCCESS); btAssert(errNum==CL_SUCCESS);
m_clipHullHullKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,srcClip, "clipHullHullKernel",&errNum,satClipContactsProg); m_clipHullHullKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,srcClip, "clipHullHullKernel",&errNum,satClipContactsProg);
@@ -120,8 +131,7 @@ m_totalContactsOut(m_context, m_queue)
if (1) if (1)
{ {
const char* srcBvh = bvhTraversalKernelCL; const char* srcBvh = bvhTraversalKernelCL;
cl_program bvhTraversalProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,srcBvh,&errNum,"","opencl/gpu_sat/kernels/bvhTraversal.cl",true); cl_program bvhTraversalProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,srcBvh,&errNum,"","opencl/gpu_sat/kernels/bvhTraversal.cl");
//cl_program bvhTraversalProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,0,&errNum,"","opencl/gpu_sat/kernels/bvhTraversal.cl", true);
btAssert(errNum==CL_SUCCESS); btAssert(errNum==CL_SUCCESS);
m_bvhTraversalKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,srcBvh, "bvhTraversalKernel",&errNum,bvhTraversalProg,""); m_bvhTraversalKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,srcBvh, "bvhTraversalKernel",&errNum,bvhTraversalProg,"");
@@ -131,7 +141,7 @@ m_totalContactsOut(m_context, m_queue)
{ {
const char* primitiveContactsSrc = primitiveContactsKernelsCL; const char* primitiveContactsSrc = primitiveContactsKernelsCL;
cl_program primitiveContactsProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,primitiveContactsSrc,&errNum,"","opencl/gpu_sat/kernels/primitiveContacts.cl",true); cl_program primitiveContactsProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,primitiveContactsSrc,&errNum,"","opencl/gpu_sat/kernels/primitiveContacts.cl");
btAssert(errNum==CL_SUCCESS); btAssert(errNum==CL_SUCCESS);
m_primitiveContactsKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,primitiveContactsSrc, "primitiveContactsKernel",&errNum,primitiveContactsProg,""); m_primitiveContactsKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,primitiveContactsSrc, "primitiveContactsKernel",&errNum,primitiveContactsProg,"");
@@ -1132,12 +1142,11 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
launcher.launch1D( num); launcher.launch1D( num);
clFinish(m_queue); clFinish(m_queue);
btAlignedObjectArray<btVector3> cpuCompoundSepNormals; // btAlignedObjectArray<btVector3> cpuCompoundSepNormals;
concaveSepNormals.copyToHost(cpuCompoundSepNormals); // concaveSepNormals.copyToHost(cpuCompoundSepNormals);
btAlignedObjectArray<btInt4> cpuConcavePairs; // btAlignedObjectArray<btInt4> cpuConcavePairs;
triangleConvexPairsOut.copyToHost(cpuConcavePairs); // triangleConvexPairsOut.copyToHost(cpuConcavePairs);
printf("!\n");
} }
} }
@@ -1345,6 +1354,9 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
launcher.launch1D( num); launcher.launch1D( num);
clFinish(m_queue); clFinish(m_queue);
nContacts = m_totalContactsOut.at(0); nContacts = m_totalContactsOut.at(0);
contactOut->resize(nContacts);
btAlignedObjectArray<btContact4> cpuContacts;
contactOut->copyToHost(cpuContacts);
// printf("nContacts after = %d\n", nContacts); // printf("nContacts after = %d\n", nContacts);
} }

15
opencl/gpu_sat/host/btConvexUtility.cpp
View File

@@ -357,12 +357,18 @@ void btConvexUtility::initialize()
edge.normalize(); edge.normalize();
bool found = false; bool found = false;
btVector3 diff,diff2;
for (int p=0;p<m_uniqueEdges.size();p++) for (int p=0;p<m_uniqueEdges.size();p++)
{ {
diff = m_uniqueEdges[p]-edge;
diff2 = m_uniqueEdges[p]+edge;
if (IsAlmostZero(m_uniqueEdges[p]-edge) || // if ((diff.length2()==0.f) ||
IsAlmostZero(m_uniqueEdges[p]+edge)) // (diff2.length2()==0.f))
if (IsAlmostZero(diff) ||
IsAlmostZero(diff2))
{ {
found = true; found = true;
break; break;
@@ -376,8 +382,9 @@ void btConvexUtility::initialize()
if (edptr) if (edptr)
{ {
btAssert(edptr->m_face0>=0); //TBD: figure out why I added this assert
btAssert(edptr->m_face1<0); // btAssert(edptr->m_face0>=0);
// btAssert(edptr->m_face1<0);
edptr->m_face1 = i; edptr->m_face1 = i;
} else } else
{ {

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

@@ -1222,13 +1222,10 @@ __kernel void findConcaveSeparatingAxisKernel( __global int4* concavePairs,
posA.w = 0.f; posA.w = 0.f;
float4 posB = rigidBodies[bodyIndexB].m_pos; float4 posB = rigidBodies[bodyIndexB].m_pos;
posB.w = 0.f; posB.w = 0.f;
float4 c0local = convexPolyhedronA.m_localCenter;
float4 ornA = rigidBodies[bodyIndexA].m_quat; float4 ornA = rigidBodies[bodyIndexA].m_quat;
float4 c0 = transform(&c0local, &posA, &ornA);
float4 c1local = convexShapes[shapeIndexB].m_localCenter;
float4 ornB =rigidBodies[bodyIndexB].m_quat; float4 ornB =rigidBodies[bodyIndexB].m_quat;
float4 c1 = transform(&c1local,&posB,&ornB);
const float4 DeltaC2 = c0 - c1;
@@ -1250,6 +1247,12 @@ __kernel void findConcaveSeparatingAxisKernel( __global int4* concavePairs,
} }
////////////////// //////////////////
float4 c0local = convexPolyhedronA.m_localCenter;
float4 c0 = transform(&c0local, &posA, &ornA);
float4 c1local = convexShapes[shapeIndexB].m_localCenter;
float4 c1 = transform(&c1local,&posB,&ornB);
const float4 DeltaC2 = c0 - c1;
bool sepA = findSeparatingAxisLocalA( &convexPolyhedronA, &convexShapes[shapeIndexB], bool sepA = findSeparatingAxisLocalA( &convexPolyhedronA, &convexShapes[shapeIndexB],
posA,ornA, posA,ornA,

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

@@ -1404,13 +1404,9 @@ __kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
posA.w = 0.f; posA.w = 0.f;
float4 posB = rigidBodies[bodyIndexB].m_pos; float4 posB = rigidBodies[bodyIndexB].m_pos;
posB.w = 0.f; posB.w = 0.f;
float4 c0local = convexPolyhedronA.m_localCenter;
float4 ornA = rigidBodies[bodyIndexA].m_quat; float4 ornA = rigidBodies[bodyIndexA].m_quat;
float4 c0 = transform(&c0local, &posA, &ornA);
float4 c1local = convexShapes[shapeIndexB].m_localCenter;
float4 ornB =rigidBodies[bodyIndexB].m_quat; float4 ornB =rigidBodies[bodyIndexB].m_quat;
float4 c1 = transform(&c1local,&posB,&ornB);
const float4 DeltaC2 = c0 - c1;
float4 sepAxis = separatingNormals[i]; float4 sepAxis = separatingNormals[i];
@@ -1419,16 +1415,17 @@ __kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
{ {
/////////////////// ///////////////////
///compound shape support ///compound shape support
int compoundChild = concavePairsIn[pairIndex].w;
int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+compoundChild; int childShapeIndexB = concavePairsIn[pairIndex].w;
int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex; int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;
shapeIndexB = collidables[childColIndexB].m_shapeIndex;
float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition; float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;
float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation; float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;
float4 newPosB = transform(&childPosB,&posB,&ornB); float4 newPosB = transform(&childPosB,&posB,&ornB);
float4 newOrnB = qtMul(ornB,childOrnB); float4 newOrnB = qtMul(ornB,childOrnB);
posB = newPosB; posB = newPosB;
ornB = newOrnB; ornB = newOrnB;
shapeIndexB = collidables[childColIndexB].m_shapeIndex;
} }
//////////////////////////////////////// ////////////////////////////////////////

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

@@ -1406,13 +1406,9 @@ static const char* satClipKernelsCL= \
" posA.w = 0.f;\n" " posA.w = 0.f;\n"
" float4 posB = rigidBodies[bodyIndexB].m_pos;\n" " float4 posB = rigidBodies[bodyIndexB].m_pos;\n"
" posB.w = 0.f;\n" " posB.w = 0.f;\n"
" float4 c0local = convexPolyhedronA.m_localCenter;\n"
" float4 ornA = rigidBodies[bodyIndexA].m_quat;\n" " float4 ornA = rigidBodies[bodyIndexA].m_quat;\n"
" float4 c0 = transform(&c0local, &posA, &ornA);\n"
" float4 c1local = convexShapes[shapeIndexB].m_localCenter;\n"
" float4 ornB =rigidBodies[bodyIndexB].m_quat;\n" " float4 ornB =rigidBodies[bodyIndexB].m_quat;\n"
" float4 c1 = transform(&c1local,&posB,&ornB);\n" "\n"
" const float4 DeltaC2 = c0 - c1;\n"
"\n" "\n"
" float4 sepAxis = separatingNormals[i];\n" " float4 sepAxis = separatingNormals[i];\n"
" \n" " \n"
@@ -1421,16 +1417,17 @@ static const char* satClipKernelsCL= \
" {\n" " {\n"
" ///////////////////\n" " ///////////////////\n"
" ///compound shape support\n" " ///compound shape support\n"
" int compoundChild = concavePairsIn[pairIndex].w;\n" " \n"
" int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+compoundChild;\n" " int 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"
" float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;\n" " float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;\n"
" float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;\n" " float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;\n"
" float4 newPosB = transform(&childPosB,&posB,&ornB);\n" " float4 newPosB = transform(&childPosB,&posB,&ornB);\n"
" float4 newOrnB = qtMul(ornB,childOrnB);\n" " float4 newOrnB = qtMul(ornB,childOrnB);\n"
" posB = newPosB;\n" " posB = newPosB;\n"
" ornB = newOrnB;\n" " ornB = newOrnB;\n"
" shapeIndexB = collidables[childColIndexB].m_shapeIndex;\n" " \n"
" }\n" " }\n"
" \n" " \n"
" ////////////////////////////////////////\n" " ////////////////////////////////////////\n"

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

@@ -1224,13 +1224,10 @@ static const char* satKernelsCL= \
" posA.w = 0.f;\n" " posA.w = 0.f;\n"
" float4 posB = rigidBodies[bodyIndexB].m_pos;\n" " float4 posB = rigidBodies[bodyIndexB].m_pos;\n"
" posB.w = 0.f;\n" " posB.w = 0.f;\n"
" float4 c0local = convexPolyhedronA.m_localCenter;\n" "\n"
" float4 ornA = rigidBodies[bodyIndexA].m_quat;\n" " float4 ornA = rigidBodies[bodyIndexA].m_quat;\n"
" float4 c0 = transform(&c0local, &posA, &ornA);\n"
" float4 c1local = convexShapes[shapeIndexB].m_localCenter;\n"
" float4 ornB =rigidBodies[bodyIndexB].m_quat;\n" " float4 ornB =rigidBodies[bodyIndexB].m_quat;\n"
" float4 c1 = transform(&c1local,&posB,&ornB);\n" "\n"
" const float4 DeltaC2 = c0 - c1;\n"
" \n" " \n"
"\n" "\n"
"\n" "\n"
@@ -1252,6 +1249,12 @@ static const char* satKernelsCL= \
" }\n" " }\n"
" //////////////////\n" " //////////////////\n"
"\n" "\n"
" float4 c0local = convexPolyhedronA.m_localCenter;\n"
" float4 c0 = transform(&c0local, &posA, &ornA);\n"
" float4 c1local = convexShapes[shapeIndexB].m_localCenter;\n"
" float4 c1 = transform(&c1local,&posB,&ornB);\n"
" const float4 DeltaC2 = c0 - c1;\n"
"\n"
"\n" "\n"
" bool sepA = findSeparatingAxisLocalA( &convexPolyhedronA, &convexShapes[shapeIndexB],\n" " bool sepA = findSeparatingAxisLocalA( &convexPolyhedronA, &convexShapes[shapeIndexB],\n"
" posA,ornA,\n" " posA,ornA,\n"