diff --git a/demo/gpudemo/GpuDemo.h b/demo/gpudemo/GpuDemo.h index c8ef1303c..4b34bf868 100644 --- a/demo/gpudemo/GpuDemo.h +++ b/demo/gpudemo/GpuDemo.h @@ -38,9 +38,9 @@ public: :useOpenCL(true), preferredOpenCLPlatformIndex(-1), preferredOpenCLDeviceIndex(-1), - arraySizeX(1), - arraySizeY(2), - arraySizeZ(1), + arraySizeX(10), + arraySizeY(30), + arraySizeZ(10), m_useConcaveMesh(false), gapX(14.3), gapY(14.0), diff --git a/demo/gpudemo/main_opengl3core.cpp b/demo/gpudemo/main_opengl3core.cpp index 313fe052c..b22f61fde 100644 --- a/demo/gpudemo/main_opengl3core.cpp +++ b/demo/gpudemo/main_opengl3core.cpp @@ -65,7 +65,7 @@ btAlignedObjectArray demoNames; int selectedDemo = 0; GpuDemo::CreateFunc* allDemos[]= { - ConcaveCompound2Scene::MyCreateFunc, +// ConcaveCompound2Scene::MyCreateFunc, ConcaveCompoundScene::MyCreateFunc, diff --git a/demo/gpudemo/rigidbody/ConcaveScene.cpp b/demo/gpudemo/rigidbody/ConcaveScene.cpp index 34ebc95ec..653276307 100644 --- a/demo/gpudemo/rigidbody/ConcaveScene.cpp +++ b/demo/gpudemo/rigidbody/ConcaveScene.cpp @@ -160,12 +160,12 @@ void ConcaveScene::createConcaveMesh(const ConstructionInfo& ci) objLoader* objData = new objLoader(); //char* fileName = "data/slopedPlane100.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/sponza_closed.obj";//"plane.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"; btVector3 shift(0,0,0);//0,230,80);//150,-100,-120); @@ -324,7 +324,7 @@ void ConcaveScene::createDynamicObjects(const ConstructionInfo& ci) float mass = 1; //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); btVector4 color = colors[curColor]; @@ -359,8 +359,8 @@ void ConcaveCompound2Scene::createDynamicObjects(const ConstructionInfo& ci) { objLoader* objData = new objLoader(); - //char* fileName = "data/teddy2_VHACD_CHs.obj"; - char* fileName = "data/cube_offset.obj"; + char* fileName = "data/teddy2_VHACD_CHs.obj"; + //char* fileName = "data/cube_offset.obj"; 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); 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; //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(btVector3(0,0,1),1.3); + btQuaternion orn(btVector3(0,0,1),1.8); btVector4 color = colors[curColor]; curColor++; @@ -665,7 +665,7 @@ btVector3 childPositions[3] = { { 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(btVector3(1,0,0),0.7); diff --git a/opencl/gpu_sat/host/ConvexHullContact.cpp b/opencl/gpu_sat/host/ConvexHullContact.cpp index efc9e0598..a60af2534 100644 --- a/opencl/gpu_sat/host/ConvexHullContact.cpp +++ b/opencl/gpu_sat/host/ConvexHullContact.cpp @@ -53,12 +53,17 @@ m_totalContactsOut(m_context, m_queue) m_totalContactsOut.push_back(0); cl_int errNum=0; - bool disableKernelCaching = true; if (1) { 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); m_findSeparatingAxisKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,src, "findSeparatingAxisKernel",&errNum,satProg ); @@ -80,7 +85,13 @@ m_totalContactsOut(m_context, m_queue) if (1) { 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); m_clipHullHullKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,srcClip, "clipHullHullKernel",&errNum,satClipContactsProg); @@ -120,8 +131,7 @@ m_totalContactsOut(m_context, m_queue) if (1) { 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,0,&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"); btAssert(errNum==CL_SUCCESS); 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; - 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); m_primitiveContactsKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,primitiveContactsSrc, "primitiveContactsKernel",&errNum,primitiveContactsProg,""); @@ -1132,12 +1142,11 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray cpuCompoundSepNormals; - concaveSepNormals.copyToHost(cpuCompoundSepNormals); - btAlignedObjectArray cpuConcavePairs; - triangleConvexPairsOut.copyToHost(cpuConcavePairs); +// btAlignedObjectArray cpuCompoundSepNormals; + // concaveSepNormals.copyToHost(cpuCompoundSepNormals); + // btAlignedObjectArray cpuConcavePairs; + // triangleConvexPairsOut.copyToHost(cpuConcavePairs); - printf("!\n"); } } @@ -1345,7 +1354,10 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArrayresize(nContacts); + btAlignedObjectArray cpuContacts; + contactOut->copyToHost(cpuContacts); +// printf("nContacts after = %d\n", nContacts); } diff --git a/opencl/gpu_sat/host/btConvexUtility.cpp b/opencl/gpu_sat/host/btConvexUtility.cpp index 7664ac348..e54942e1a 100644 --- a/opencl/gpu_sat/host/btConvexUtility.cpp +++ b/opencl/gpu_sat/host/btConvexUtility.cpp @@ -357,12 +357,18 @@ void btConvexUtility::initialize() edge.normalize(); bool found = false; + btVector3 diff,diff2; for (int p=0;pm_face0>=0); - btAssert(edptr->m_face1<0); + //TBD: figure out why I added this assert +// btAssert(edptr->m_face0>=0); + // btAssert(edptr->m_face1<0); edptr->m_face1 = i; } else { diff --git a/opencl/gpu_sat/kernels/sat.cl b/opencl/gpu_sat/kernels/sat.cl index e06fc85d6..e8e9e68a1 100644 --- a/opencl/gpu_sat/kernels/sat.cl +++ b/opencl/gpu_sat/kernels/sat.cl @@ -1222,13 +1222,10 @@ __kernel void findConcaveSeparatingAxisKernel( __global int4* concavePairs, posA.w = 0.f; float4 posB = rigidBodies[bodyIndexB].m_pos; posB.w = 0.f; - float4 c0local = convexPolyhedronA.m_localCenter; + 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 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], posA,ornA, diff --git a/opencl/gpu_sat/kernels/satClipHullContacts.cl b/opencl/gpu_sat/kernels/satClipHullContacts.cl index 96fbdfa25..78f83719f 100644 --- a/opencl/gpu_sat/kernels/satClipHullContacts.cl +++ b/opencl/gpu_sat/kernels/satClipHullContacts.cl @@ -1404,13 +1404,9 @@ __kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn, posA.w = 0.f; float4 posB = rigidBodies[bodyIndexB].m_pos; posB.w = 0.f; - float4 c0local = convexPolyhedronA.m_localCenter; 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 c1 = transform(&c1local,&posB,&ornB); - const float4 DeltaC2 = c0 - c1; + float4 sepAxis = separatingNormals[i]; @@ -1419,16 +1415,17 @@ __kernel void clipHullHullConcaveConvexKernel( __global int4* concavePairsIn, { /////////////////// ///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; + shapeIndexB = collidables[childColIndexB].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; + } //////////////////////////////////////// diff --git a/opencl/gpu_sat/kernels/satClipHullContacts.h b/opencl/gpu_sat/kernels/satClipHullContacts.h index 4dc53cf8d..49030b6c7 100644 --- a/opencl/gpu_sat/kernels/satClipHullContacts.h +++ b/opencl/gpu_sat/kernels/satClipHullContacts.h @@ -1406,13 +1406,9 @@ static const char* satClipKernelsCL= \ " posA.w = 0.f;\n" " float4 posB = rigidBodies[bodyIndexB].m_pos;\n" " posB.w = 0.f;\n" -" float4 c0local = convexPolyhedronA.m_localCenter;\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 c1 = transform(&c1local,&posB,&ornB);\n" -" const float4 DeltaC2 = c0 - c1;\n" +"\n" "\n" " float4 sepAxis = separatingNormals[i];\n" " \n" @@ -1421,16 +1417,17 @@ static const char* satClipKernelsCL= \ " {\n" " ///////////////////\n" " ///compound shape support\n" -" int compoundChild = concavePairsIn[pairIndex].w;\n" -" int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+compoundChild;\n" +" \n" +" int childShapeIndexB = concavePairsIn[pairIndex].w;\n" " int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;\n" +" shapeIndexB = collidables[childColIndexB].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" diff --git a/opencl/gpu_sat/kernels/satKernels.h b/opencl/gpu_sat/kernels/satKernels.h index 79bf3c5cd..5e9f8b4f6 100644 --- a/opencl/gpu_sat/kernels/satKernels.h +++ b/opencl/gpu_sat/kernels/satKernels.h @@ -1224,13 +1224,10 @@ static const char* satKernelsCL= \ " posA.w = 0.f;\n" " float4 posB = rigidBodies[bodyIndexB].m_pos;\n" " posB.w = 0.f;\n" -" float4 c0local = convexPolyhedronA.m_localCenter;\n" +"\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 c1 = transform(&c1local,&posB,&ornB);\n" -" const float4 DeltaC2 = c0 - c1;\n" +"\n" " \n" "\n" "\n" @@ -1252,6 +1249,12 @@ static const char* satKernelsCL= \ " }\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" " bool sepA = findSeparatingAxisLocalA( &convexPolyhedronA, &convexShapes[shapeIndexB],\n" " posA,ornA,\n"