From b4f9416cdf240a8d71084bfb3012594fa46f1fe0 Mon Sep 17 00:00:00 2001 From: erwin coumans Date: Wed, 20 Mar 2013 23:37:34 -0700 Subject: [PATCH] add support for BVH acceleration for concave trianglemesh collision against convex hulls bugfix/improvement in batching --- btgui/OpenGLWindow/GLInstancingRenderer.cpp | 2 +- data/plane100.obj | 12 + demo/gpudemo/GpuDemo.h | 12 +- demo/gpudemo/main_opengl3core.cpp | 2 +- demo/gpudemo/rigidbody/ConcaveScene.cpp | 19 +- demo/gpudemo/rigidbody/GpuRigidBodyDemo.cpp | 8 +- opencl/gpu_rigidbody/host/Solver.cpp | 14 +- opencl/gpu_rigidbody/host/Solver.h | 2 +- opencl/gpu_rigidbody/host/btConfig.h | 5 +- .../host/btGpuBatchingPgsSolver.cpp | 126 ++--- .../host/btGpuBatchingPgsSolver.h | 2 +- .../gpu_rigidbody/host/btGpuNarrowPhase.cpp | 36 +- .../gpu_rigidbody/kernels/batchingKernels.cl | 15 +- .../gpu_rigidbody/kernels/batchingKernels.h | 15 +- opencl/gpu_sat/host/ConvexHullContact.cpp | 225 ++++----- opencl/gpu_sat/host/ConvexHullContact.h | 4 + opencl/gpu_sat/kernels/sat.cl | 473 +++++++++--------- opencl/gpu_sat/kernels/satClipHullContacts.cl | 3 + opencl/gpu_sat/kernels/satClipHullContacts.h | 3 + opencl/gpu_sat/kernels/satKernels.h | 472 ++++++++--------- 20 files changed, 760 insertions(+), 690 deletions(-) create mode 100644 data/plane100.obj diff --git a/btgui/OpenGLWindow/GLInstancingRenderer.cpp b/btgui/OpenGLWindow/GLInstancingRenderer.cpp index f99feedf0..d87d56738 100644 --- a/btgui/OpenGLWindow/GLInstancingRenderer.cpp +++ b/btgui/OpenGLWindow/GLInstancingRenderer.cpp @@ -109,7 +109,7 @@ struct InternalDataRenderer : public GLInstanceRendererInternalData m_cameraTargetPosition(btVector3(15,2,-24)), m_cameraDistance(150), m_cameraUp(0,1,0), - m_azi(135.f), + m_azi(100.f),//135.f), m_ele(25.f), m_mouseInitialized(false) { diff --git a/data/plane100.obj b/data/plane100.obj new file mode 100644 index 000000000..ae4a6c281 --- /dev/null +++ b/data/plane100.obj @@ -0,0 +1,12 @@ +# Blender v2.66 (sub 1) OBJ File: '' +# www.blender.org +mtllib plane.mtl +o Plane +v 100.000000 0.000000 -100.000000 +v 100.000000 0.000000 100.000000 +v -100.000000 0.000000 100.000000 +v -100.000000 0.000000 -100.000000 +usemtl Material +s off +f 1 2 3 +f 1 3 4 diff --git a/demo/gpudemo/GpuDemo.h b/demo/gpudemo/GpuDemo.h index 4c07e3d07..1b3ba41db 100644 --- a/demo/gpudemo/GpuDemo.h +++ b/demo/gpudemo/GpuDemo.h @@ -34,13 +34,13 @@ public: :useOpenCL(true), preferredOpenCLPlatformIndex(-1), preferredOpenCLDeviceIndex(-1), - arraySizeX(5), - arraySizeY(5 ), - arraySizeZ(5), + arraySizeX(30), + arraySizeY(10 ), + arraySizeZ(30), m_useConcaveMesh(false), - gapX(4.3), - gapY(2.0), - gapZ(4.3), + gapX(6.3), + gapY(12.0), + gapZ(6.3), m_instancingRenderer(0), m_window(0) { diff --git a/demo/gpudemo/main_opengl3core.cpp b/demo/gpudemo/main_opengl3core.cpp index 1c51f8357..a3e64a6d2 100644 --- a/demo/gpudemo/main_opengl3core.cpp +++ b/demo/gpudemo/main_opengl3core.cpp @@ -617,7 +617,7 @@ int main(int argc, char* argv[]) if (!gPause) { - BT_PROFILE("simulate"); + BT_PROFILE("clientMoveAndDisplay"); demo->clientMoveAndDisplay(); } diff --git a/demo/gpudemo/rigidbody/ConcaveScene.cpp b/demo/gpudemo/rigidbody/ConcaveScene.cpp index 3c99d6b3d..f846f94b5 100644 --- a/demo/gpudemo/rigidbody/ConcaveScene.cpp +++ b/demo/gpudemo/rigidbody/ConcaveScene.cpp @@ -166,9 +166,9 @@ GraphicsShape* createGraphicsShapeFromWavefrontObj(objLoader* obj) void ConcaveScene::setupScene(const ConstructionInfo& ci) { objLoader* objData = new objLoader(); - //char* fileName = "data/plane.obj"; - char* fileName = "data/teddy.obj";//"plane.obj"; - //char* fileName = "data/sponza_closed.obj";//"plane.obj"; + //char* fileName = "data/plane100.obj"; + //char* fileName = "data/teddy.obj";//"plane.obj"; + char* fileName = "data/sponza_closed.obj";//"plane.obj"; FILE* f = 0; @@ -200,16 +200,16 @@ void ConcaveScene::setupScene(const ConstructionInfo& ci) { GraphicsShape* shape = createGraphicsShapeFromWavefrontObj(objData); - btVector4 scaling(1,1,1,1); + btVector4 scaling(4,4,4,1); btAlignedObjectArray verts; for (int i=0;im_numvertices;i++) { btVector3 vtx = (btVector3&)shape->m_vertices->at(i).xyzw; - verts.push_back(vtx); + verts.push_back(vtx*scaling); } - int colIndex = m_data->m_np->registerConcaveMesh(&verts,shape->m_indices,scaling); + int colIndex = m_data->m_np->registerConcaveMesh(&verts,shape->m_indices,btVector3(1,1,1)); { int strideInBytes = 9*sizeof(float); @@ -222,7 +222,7 @@ void ConcaveScene::setupScene(const ConstructionInfo& ci) int shapeId = ci.m_instancingRenderer->registerShape(&shape->m_vertices->at(0).xyzw[0], shape->m_numvertices, &shape->m_indices->at(0), shape->m_numIndices); btQuaternion orn(0,0,0,1); - btVector4 color(0,1,0,1.f);//0.5); + btVector4 color(0,0,1,1.f);//0.5);//1.f { @@ -262,7 +262,8 @@ void ConcaveScene::setupScene(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(-2*ci.gapX+i*ci.gapX,25+j*ci.gapY,-2*ci.gapZ+k*ci.gapZ); + btVector3 position(-(ci.arraySizeX/2)*ci.gapX+i*ci.gapX,5+j*ci.gapY,-(ci.arraySizeZ/2)*ci.gapZ+k*ci.gapZ); btQuaternion orn(1,0,0,0); btVector4 color(0,1,0,1); @@ -278,6 +279,6 @@ void ConcaveScene::setupScene(const ConstructionInfo& ci) float camPos[4]={0,0,0,0};//65.5,4.5,65.5,0}; //float camPos[4]={1,12.5,1.5,0}; m_instancingRenderer->setCameraTargetPosition(camPos); - m_instancingRenderer->setCameraDistance(50); + m_instancingRenderer->setCameraDistance(120); } \ No newline at end of file diff --git a/demo/gpudemo/rigidbody/GpuRigidBodyDemo.cpp b/demo/gpudemo/rigidbody/GpuRigidBodyDemo.cpp index 30fb2090a..076cf73f5 100644 --- a/demo/gpudemo/rigidbody/GpuRigidBodyDemo.cpp +++ b/demo/gpudemo/rigidbody/GpuRigidBodyDemo.cpp @@ -157,6 +157,7 @@ void GpuRigidBodyDemo::clientMoveAndDisplay() btVector4* positions = 0; if (animate && numObjects) { + BT_PROFILE("gl2cl"); GLuint vbo = m_instancingRenderer->getInternalData()->m_vbo; int arraySizeInBytes = numObjects * (3)*sizeof(btVector4); glBindBuffer(GL_ARRAY_BUFFER, vbo); @@ -172,10 +173,14 @@ void GpuRigidBodyDemo::clientMoveAndDisplay() } } - m_data->m_rigidBodyPipeline->stepSimulation(1./60.f); + { + BT_PROFILE("stepSimulation"); + m_data->m_rigidBodyPipeline->stepSimulation(1./60.f); + } if (numObjects) { + BT_PROFILE("cl2gl_convert"); int ciErrNum = 0; cl_mem bodies = m_data->m_rigidBodyPipeline->getBodyBuffer(); btLauncherCL launch(m_clData->m_clQueue,m_data->m_copyTransformsToVBOKernel); @@ -188,6 +193,7 @@ void GpuRigidBodyDemo::clientMoveAndDisplay() if (animate && numObjects) { + BT_PROFILE("cl2gl_upload"); GLint err = glGetError(); assert(err==GL_NO_ERROR); m_data->m_instancePosOrnColor->copyToHostPointer(positions,3*numObjects,0); diff --git a/opencl/gpu_rigidbody/host/Solver.cpp b/opencl/gpu_rigidbody/host/Solver.cpp index 4ebb1d441..221eaf18e 100644 --- a/opencl/gpu_rigidbody/host/Solver.cpp +++ b/opencl/gpu_rigidbody/host/Solver.cpp @@ -818,9 +818,10 @@ void Solver::sortContacts( const btOpenCLArray* bodyBuf, */ -void Solver::batchContacts( btOpenCLArray* contacts, int nContacts, btOpenCLArray* nNative, btOpenCLArray* offsetsNative, int staticIdx ) +void Solver::batchContacts( btOpenCLArray* contacts, int nContacts, btOpenCLArray* nNative, btOpenCLArray* offsetsNative, int staticIdx ) { - + + int numWorkItems = 64*N_SPLIT*N_SPLIT; { BT_PROFILE("batch generation"); @@ -829,7 +830,7 @@ void Solver::batchContacts( btOpenCLArray* contacts, int nContacts, cdata.y = 0; cdata.z = staticIdx; - int numWorkItems = 64*N_SPLIT*N_SPLIT; + #ifdef BATCH_DEBUG SolverDebugInfo* debugInfo = new SolverDebugInfo[numWorkItems]; adl::btOpenCLArray gpuDebugInfo(data->m_device,numWorkItems); @@ -837,12 +838,14 @@ void Solver::batchContacts( btOpenCLArray* contacts, int nContacts, gpuDebugInfo.write(debugInfo,numWorkItems); #endif + + btBufferInfoCL bInfo[] = { btBufferInfoCL( contacts->getBufferCL() ), btBufferInfoCL( m_contactBuffer->getBufferCL() ), btBufferInfoCL( nNative->getBufferCL() ), - btBufferInfoCL( offsetsNative->getBufferCL() ) + btBufferInfoCL( offsetsNative->getBufferCL() ), #ifdef BATCH_DEBUG , btBufferInfoCL(&gpuDebugInfo) #endif @@ -899,7 +902,8 @@ void Solver::batchContacts( btOpenCLArray* contacts, int nContacts, btAssert(m_contactBuffer->size()==nContacts); //contacts->copyFromOpenCLArray( *m_contactBuffer); //clFinish(m_queue);//needed? - + + } diff --git a/opencl/gpu_rigidbody/host/Solver.h b/opencl/gpu_rigidbody/host/Solver.h index b027dfc3b..ca6a8d4ad 100644 --- a/opencl/gpu_rigidbody/host/Solver.h +++ b/opencl/gpu_rigidbody/host/Solver.h @@ -132,7 +132,7 @@ class Solver : public SolverBase btOpenCLArray* contactsIn, btOpenCLArray* contactCOut, void* additionalData, int nContacts, const ConstraintCfg& cfg ); - void batchContacts( btOpenCLArray* contacts, int nContacts, btOpenCLArray* n, btOpenCLArray* offsets, int staticIdx ); + void batchContacts( btOpenCLArray* contacts, int nContacts, btOpenCLArray* n, btOpenCLArray* offsets, int staticIdx ); }; diff --git a/opencl/gpu_rigidbody/host/btConfig.h b/opencl/gpu_rigidbody/host/btConfig.h index 71b1f1b6c..17cc98990 100644 --- a/opencl/gpu_rigidbody/host/btConfig.h +++ b/opencl/gpu_rigidbody/host/btConfig.h @@ -14,6 +14,8 @@ struct btConfig int m_maxConvexUniqueEdges; int m_maxCompoundChildShapes; + + int m_maxTriConvexPairCapacity; btConfig() :m_maxConvexBodies(128*1024), @@ -23,7 +25,8 @@ struct btConfig m_maxConvexVertices(8192), m_maxConvexIndices(8192), m_maxConvexUniqueEdges(8192), - m_maxCompoundChildShapes(8192)//?? + m_maxCompoundChildShapes(8192), + m_maxTriConvexPairCapacity(256*1024) { m_maxBroadphasePairs = 16*m_maxConvexBodies; } diff --git a/opencl/gpu_rigidbody/host/btGpuBatchingPgsSolver.cpp b/opencl/gpu_rigidbody/host/btGpuBatchingPgsSolver.cpp index 3e7b92368..170bc40df 100644 --- a/opencl/gpu_rigidbody/host/btGpuBatchingPgsSolver.cpp +++ b/opencl/gpu_rigidbody/host/btGpuBatchingPgsSolver.cpp @@ -202,7 +202,7 @@ btGpuBatchingPgsSolver::~btGpuBatchingPgsSolver() struct btConstraintCfg { - btConstraintCfg( float dt = 0.f ): m_positionDrift( 0.005f ), m_positionConstraintCoeff( 0.2f ), m_dt(dt), m_staticIdx(-1) {} + btConstraintCfg( float dt = 0.f ): m_positionDrift( 0.005f ), m_positionConstraintCoeff( 0.2f ), m_dt(dt), m_staticIdx(0) {} float m_positionDrift; float m_positionConstraintCoeff; @@ -403,8 +403,8 @@ void btGpuBatchingPgsSolver::solveContacts(int numBodies, cl_mem bodyBuf, cl_mem float dt=1./60.; btConstraintCfg csCfg( dt ); csCfg.m_enableParallelSolve = true; - csCfg.m_averageExtent = 0.2f;//@TODO m_averageObjExtent; - csCfg.m_staticIdx = -1;//m_static0Index;//m_planeBodyIndex; + csCfg.m_averageExtent = .2f;//@TODO m_averageObjExtent; + csCfg.m_staticIdx = 0;//m_static0Index;//m_planeBodyIndex; btOpenCLArray* contactsIn = m_data->m_pBufContactOutGPU; btOpenCLArray* bodyBuf = m_data->m_bodyBufferGPU; @@ -572,64 +572,65 @@ void btGpuBatchingPgsSolver::solveContacts(int numBodies, cl_mem bodyBuf, cl_mem bool compareGPU = false; if (nContacts) { - if (gpuBatchContacts) - { - maxNumBatches=250;//for now - BT_PROFILE("gpu batchContacts"); - m_data->m_solverGPU->batchContacts( (btOpenCLArray*)contactNative, nContacts, m_data->m_solverGPU->m_numConstraints, m_data->m_solverGPU->m_offsets, csCfg.m_staticIdx ); - } else - { - BT_PROFILE("cpu batchContacts"); - btAlignedObjectArray cpuContacts; - btOpenCLArray* contactsIn = m_data->m_pBufContactOutGPU; - contactsIn->copyToHost(cpuContacts); + if (gpuBatchContacts) + { + BT_PROFILE("gpu batchContacts"); + maxNumBatches = 50; + m_data->m_solverGPU->batchContacts( (btOpenCLArray*)contactNative, nContacts, m_data->m_solverGPU->m_numConstraints, m_data->m_solverGPU->m_offsets, csCfg.m_staticIdx ); + } else + { + BT_PROFILE("cpu batchContacts"); + btAlignedObjectArray cpuContacts; + btOpenCLArray* contactsIn = m_data->m_pBufContactOutGPU; + contactsIn->copyToHost(cpuContacts); - btOpenCLArray* countsNative = m_data->m_solverGPU->m_numConstraints; - btOpenCLArray* offsetsNative = m_data->m_solverGPU->m_offsets; + btOpenCLArray* countsNative = m_data->m_solverGPU->m_numConstraints; + btOpenCLArray* offsetsNative = m_data->m_solverGPU->m_offsets; - btAlignedObjectArray nNativeHost; - btAlignedObjectArray offsetsNativeHost; + btAlignedObjectArray nNativeHost; + btAlignedObjectArray offsetsNativeHost; - { - BT_PROFILE("countsNative/offsetsNative copyToHost"); - countsNative->copyToHost(nNativeHost); - offsetsNative->copyToHost(offsetsNativeHost); - } + { + BT_PROFILE("countsNative/offsetsNative copyToHost"); + countsNative->copyToHost(nNativeHost); + offsetsNative->copyToHost(offsetsNativeHost); + } - int numNonzeroGrid=0; + int numNonzeroGrid=0; - { - BT_PROFILE("batch grid"); - for(int i=0; im_queue); + clFinish(m_data->m_queue); - } - } - } - { - BT_PROFILE("m_contactBuffer->copyFromHost"); - m_data->m_solverGPU->m_contactBuffer->copyFromHost((btAlignedObjectArray&)cpuContacts); - } - // printf("maxNumBatches = %d\n", maxNumBatches); - } + } + } + } + { + BT_PROFILE("m_contactBuffer->copyFromHost"); + m_data->m_solverGPU->m_contactBuffer->copyFromHost((btAlignedObjectArray&)cpuContacts); + } + + } + } - + //printf("maxNumBatches = %d\n", maxNumBatches); if (nContacts) { @@ -704,8 +705,14 @@ btAlignedObjectArray sortData; btAlignedObjectArray old; -inline int btGpuBatchingPgsSolver::sortConstraintByBatch( btContact4* cs, int n, int simdWidth , int staticIdx) +inline int btGpuBatchingPgsSolver::sortConstraintByBatch( btContact4* cs, int n, int simdWidth , int staticIdx, int numBodies) { + btAlignedObjectArray bodyUsed; + bodyUsed.resize(numBodies); + for (int q=0;q=0)&&bodyAS!=staticIdx? aUnavailable:0;// - bUnavailable = (bodyBS>=0)&&bodyBS!=staticIdx? bUnavailable:0; + aUnavailable = !aIsStatic? aUnavailable:0;// + bUnavailable = !bIsStatic? bUnavailable:0; if( aUnavailable==0 && bUnavailable==0 ) // ok { - flg[ aIdx/32 ] |= (1<<(aIdx&31)); - flg[ bIdx/32 ] |= (1<<(bIdx&31)); + if (!!aIsStatic) + flg[ aIdx/32 ] |= (1<<(aIdx&31)); + if (!bIsStatic) + flg[ bIdx/32 ] |= (1<<(bIdx&31)); + cs[idx].getBatchIdx() = batchIdx; sortData[idx].m_key = batchIdx; sortData[idx].m_value = idx; diff --git a/opencl/gpu_rigidbody/host/btGpuBatchingPgsSolver.h b/opencl/gpu_rigidbody/host/btGpuBatchingPgsSolver.h index 4bb697100..2ebc4cf8e 100644 --- a/opencl/gpu_rigidbody/host/btGpuBatchingPgsSolver.h +++ b/opencl/gpu_rigidbody/host/btGpuBatchingPgsSolver.h @@ -15,7 +15,7 @@ protected: struct btGpuBatchingPgsSolverInternalData* m_data; void batchContacts( btOpenCLArray* contacts, int nContacts, btOpenCLArray* n, btOpenCLArray* offsets, int staticIdx ); - inline int sortConstraintByBatch( btContact4* cs, int n, int simdWidth , int staticIdx); + inline int sortConstraintByBatch( btContact4* cs, int n, int simdWidth , int staticIdx, int numBodies); void solveContactConstraint( const btOpenCLArray* bodyBuf, const btOpenCLArray* shapeBuf, btOpenCLArray* constraint, void* additionalData, int n ,int maxNumBatches, int numIterations); diff --git a/opencl/gpu_rigidbody/host/btGpuNarrowPhase.cpp b/opencl/gpu_rigidbody/host/btGpuNarrowPhase.cpp index beec30c51..306ce903e 100644 --- a/opencl/gpu_rigidbody/host/btGpuNarrowPhase.cpp +++ b/opencl/gpu_rigidbody/host/btGpuNarrowPhase.cpp @@ -63,6 +63,10 @@ struct btGpuNarrowPhaseInternalData btAlignedObjectArray* m_localShapeAABBCPU; btAlignedObjectArray m_bvhData; + btOpenCLArray* m_treeNodesGPU; + btOpenCLArray* m_subTreesGPU; + + btConfig m_config; }; @@ -137,6 +141,10 @@ m_queue(queue) m_data->m_numAcceleratedShapes = 0; m_data->m_numAcceleratedRigidBodies = 0; + m_data->m_treeNodesGPU = 0; + m_data->m_subTreesGPU = 0; + + //m_data->m_contactCGPU = new btOpenCLArray(ctx,queue,config.m_maxBroadphasePairs,false); //m_data->m_frictionCGPU = new btOpenCLArray::allocateFrictionConstraint( m_data->m_deviceCL, config.m_maxBroadphasePairs); @@ -170,6 +178,8 @@ btGpuNarrowPhase::~btGpuNarrowPhase() delete m_data->m_worldVertsA1GPU; delete m_data->m_worldVertsB2GPU; + delete m_data->m_treeNodesGPU; + delete m_data->m_subTreesGPU; delete m_data->m_convexData; delete m_data; @@ -328,6 +338,15 @@ int btGpuNarrowPhase::registerConvexHullShape(btConvexUtility* utilPtr) int btGpuNarrowPhase::registerConcaveMesh(btAlignedObjectArray* vertices, btAlignedObjectArray* indices,const float* scaling1) { + //right now we only support one single mesh, it is on the todo to merge all mesh data etc + btAssert(m_data->m_treeNodesGPU ==0); + btAssert(m_data->m_subTreesGPU ==0); + if (m_data->m_treeNodesGPU) + { + printf("error, only 1 single concave mesh supported at the moment\n"); + exit (0); + } + btVector3 scaling(scaling1[0],scaling1[1],scaling1[2]); int collidableIndex = allocateCollidable(); @@ -377,6 +396,17 @@ int btGpuNarrowPhase::registerConcaveMesh(btAlignedObjectArray* vert meshInterface->addIndexedMesh(mesh); bvh->build(meshInterface, useQuantizedAabbCompression, (btVector3&)aabb.m_min, (btVector3&)aabb.m_max); m_data->m_bvhData.push_back(bvh); + int numNodes = bvh->getQuantizedNodeArray().size(); + btOpenCLArray* treeNodesGPU = new btOpenCLArray(this->m_context,this->m_queue,numNodes); + treeNodesGPU->copyFromHost(bvh->getQuantizedNodeArray()); + + int numSubTrees = bvh->getSubtreeInfoArray().size(); + btOpenCLArray* subTreesGPU = new btOpenCLArray(this->m_context,this->m_queue,numSubTrees); + subTreesGPU->copyFromHost(bvh->getSubtreeInfoArray()); + + m_data->m_treeNodesGPU = treeNodesGPU; + m_data->m_subTreesGPU = subTreesGPU; + return collidableIndex; } @@ -412,7 +442,7 @@ int btGpuNarrowPhase::registerConcaveMeshShape(btAlignedObjectArray* { if (i%256==0) { - printf("i=%d out of %d", i,convex.m_numFaces); + //printf("i=%d out of %d", i,convex.m_numFaces); } btVector3 vert0(vertices->at(indices->at(i*3))*scaling); btVector3 vert1(vertices->at(indices->at(i*3+1))*scaling); @@ -524,7 +554,7 @@ void btGpuNarrowPhase::computeContacts(cl_mem broadphasePairs, int numBroadphase { int nContactOut = 0; - int maxTriConvexPairCapacity = 8192; + int maxTriConvexPairCapacity = m_data->m_config.m_maxTriConvexPairCapacity; btOpenCLArray triangleConvexPairs(m_context,m_queue, maxTriConvexPairCapacity); int numTriConvexPairsOut=0; @@ -552,6 +582,8 @@ void btGpuNarrowPhase::computeContacts(cl_mem broadphasePairs, int numBroadphase *m_data->m_worldVertsA1GPU, *m_data->m_worldVertsB2GPU, m_data->m_bvhData, + m_data->m_treeNodesGPU, + m_data->m_subTreesGPU, numObjects, maxTriConvexPairCapacity, triangleConvexPairs, diff --git a/opencl/gpu_rigidbody/kernels/batchingKernels.cl b/opencl/gpu_rigidbody/kernels/batchingKernels.cl index 988d8c932..5087f23b0 100644 --- a/opencl/gpu_rigidbody/kernels/batchingKernels.cl +++ b/opencl/gpu_rigidbody/kernels/batchingKernels.cl @@ -224,17 +224,20 @@ __kernel void CreateBatches( __global const Contact4* gConstraints, __global Con aAvailable = tryWrite( ldsCheckBuffer, ea ); bAvailable = tryWrite( ldsCheckBuffer, eb ); - aAvailable = (e.m_a<0)? 1: aAvailable; - bAvailable = (e.m_b<0)? 1: bAvailable; + bool aStatic = (e.m_a<0) ||(ea==m_staticIdx); + bool bStatic = (e.m_b<0) ||(eb==m_staticIdx); - aAvailable = (e.m_a==m_staticIdx)? 1: aAvailable; - bAvailable = (e.m_b==m_staticIdx)? 1: bAvailable; + aAvailable = aStatic? 1: aAvailable; + bAvailable = bStatic? 1: bAvailable; bool success = (aAvailable && bAvailable); if(success) { - writeBuf( ldsFixedBuffer, ea ); - writeBuf( ldsFixedBuffer, eb ); + + if (!aStatic) + writeBuf( ldsFixedBuffer, ea ); + if (!bStatic) + writeBuf( ldsFixedBuffer, eb ); } done = success; } diff --git a/opencl/gpu_rigidbody/kernels/batchingKernels.h b/opencl/gpu_rigidbody/kernels/batchingKernels.h index 4f278340c..c21a5fdaf 100644 --- a/opencl/gpu_rigidbody/kernels/batchingKernels.h +++ b/opencl/gpu_rigidbody/kernels/batchingKernels.h @@ -226,17 +226,20 @@ static const char* batchingKernelsCL= \ " aAvailable = tryWrite( ldsCheckBuffer, ea );\n" " bAvailable = tryWrite( ldsCheckBuffer, eb );\n" "\n" -" aAvailable = (e.m_a<0)? 1: aAvailable;\n" -" bAvailable = (e.m_b<0)? 1: bAvailable;\n" +" bool aStatic = (e.m_a<0) ||(ea==m_staticIdx);\n" +" bool bStatic = (e.m_b<0) ||(eb==m_staticIdx);\n" " \n" -" aAvailable = (e.m_a==m_staticIdx)? 1: aAvailable;\n" -" bAvailable = (e.m_b==m_staticIdx)? 1: bAvailable;\n" +" aAvailable = aStatic? 1: aAvailable;\n" +" bAvailable = bStatic? 1: bAvailable;\n" "\n" " bool success = (aAvailable && bAvailable);\n" " if(success)\n" " {\n" -" writeBuf( ldsFixedBuffer, ea );\n" -" writeBuf( ldsFixedBuffer, eb );\n" +" \n" +" if (!aStatic)\n" +" writeBuf( ldsFixedBuffer, ea );\n" +" if (!bStatic)\n" +" writeBuf( ldsFixedBuffer, eb );\n" " }\n" " done = success;\n" " }\n" diff --git a/opencl/gpu_sat/host/ConvexHullContact.cpp b/opencl/gpu_sat/host/ConvexHullContact.cpp index bcde25850..96479f519 100644 --- a/opencl/gpu_sat/host/ConvexHullContact.cpp +++ b/opencl/gpu_sat/host/ConvexHullContact.cpp @@ -23,7 +23,7 @@ subject to the following restrictions: #include "ConvexHullContact.h" #include //memcpy #include "btConvexPolyhedronCL.h" -#include "btOptimizedBvh.h" + typedef btAlignedObjectArray btVertexArray; #include "BulletCommon/btQuickprof.h" @@ -61,6 +61,8 @@ m_totalContactsOut(m_context, m_queue) m_findSeparatingAxisKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,src, "findSeparatingAxisKernel",&errNum,satProg ); + m_findConcaveSeparatingAxisKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,src, "findConcaveSeparatingAxisKernel",&errNum,satProg ); + m_findCompoundPairsKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,src, "findCompoundPairsKernel",&errNum,satProg ); @@ -130,6 +132,9 @@ GpuSatCollision::~GpuSatCollision() if (m_findSeparatingAxisKernel) clReleaseKernel(m_findSeparatingAxisKernel); + if (m_findConcaveSeparatingAxisKernel) + clReleaseKernel(m_findConcaveSeparatingAxisKernel); + if (m_findCompoundPairsKernel) clReleaseKernel(m_findCompoundPairsKernel); @@ -190,6 +195,8 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray& worldVertsA1GPU, btOpenCLArray& worldVertsB2GPU, btAlignedObjectArray& bvhData, + btOpenCLArray* treeNodesGPU, + btOpenCLArray* subTreesGPU, int numObjects, int maxTriConvexPairCapacity, btOpenCLArray& triangleConvexPairsOut, @@ -231,146 +238,112 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArraygetBufferCL(), true ), - btBufferInfoCL( bodyBuf->getBufferCL(),true), - btBufferInfoCL( gpuCollidables.getBufferCL(),true), - btBufferInfoCL( convexData.getBufferCL(),true), - btBufferInfoCL( gpuVertices.getBufferCL(),true), - btBufferInfoCL( gpuUniqueEdges.getBufferCL(),true), - btBufferInfoCL( gpuFaces.getBufferCL(),true), - btBufferInfoCL( gpuIndices.getBufferCL(),true), - btBufferInfoCL( clAabbsWS.getBufferCL(),true), - btBufferInfoCL( sepNormals.getBufferCL()), - btBufferInfoCL( hasSeparatingNormals.getBufferCL()), - btBufferInfoCL( triangleConvexPairsOut.getBufferCL()), - btBufferInfoCL( concaveSepNormals.getBufferCL()), - btBufferInfoCL( numConcavePairsOut.getBufferCL()) - }; - - btLauncherCL launcher(m_queue, m_findSeparatingAxisKernel); - launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(btBufferInfoCL) ); - launcher.setConst( nPairs ); - launcher.setConst( maxTriConvexPairCapacity); - - int num = nPairs; - launcher.launch1D( num); - clFinish(m_queue); - - numConcave = numConcavePairsOut.at(0); - if (numConcave) { - if (numConcave > maxTriConvexPairCapacity) - numConcave = maxTriConvexPairCapacity; + BT_PROFILE("findSeparatingAxisKernel"); + btBufferInfoCL bInfo[] = { + btBufferInfoCL( pairs->getBufferCL(), true ), + btBufferInfoCL( bodyBuf->getBufferCL(),true), + btBufferInfoCL( gpuCollidables.getBufferCL(),true), + btBufferInfoCL( convexData.getBufferCL(),true), + btBufferInfoCL( gpuVertices.getBufferCL(),true), + btBufferInfoCL( gpuUniqueEdges.getBufferCL(),true), + btBufferInfoCL( gpuFaces.getBufferCL(),true), + btBufferInfoCL( gpuIndices.getBufferCL(),true), + btBufferInfoCL( clAabbsWS.getBufferCL(),true), + btBufferInfoCL( sepNormals.getBufferCL()), + btBufferInfoCL( hasSeparatingNormals.getBufferCL()) + }; - triangleConvexPairsOut.resize(numConcave); - btAlignedObjectArray triangleConvexPairsOutCPU; - triangleConvexPairsOut.copyToHost(triangleConvexPairsOutCPU); - printf("-----------------------\n", numConcave); - printf("got %d concave pairs\n", numConcave); - btAssert(numConcave = triangleConvexPairsOutCPU.size()); + btLauncherCL launcher(m_queue, m_findSeparatingAxisKernel); + launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(btBufferInfoCL) ); + launcher.setConst( nPairs ); - for (int i=0;i collidablesCPU; - gpuCollidables.copyToHost(collidablesCPU); - btAlignedObjectArray bodiesCPU; - bodyBuf->copyToHost(bodiesCPU); - btAlignedObjectArray pairsCPU; + int num = nPairs; + launcher.launch1D( num); + clFinish(m_queue); + } - btAlignedObjectArray aabbsWSCPU; - clAabbsWS.copyToHost(aabbsWSCPU); - - pairs->copyToHost(pairsCPU); - MyTriangleCallback triCallback; + //now perform the tree query on GPU + { - - for (int i=0;igetQuantizedNodeArray().size() : 0; + if (numNodes) { - int bodyIndexA = pairsCPU[i].x; - int bodyIndexB = pairsCPU[i].y; - - triCallback.m_bodyIndexA = bodyIndexA; - triCallback.m_bodyIndexB = bodyIndexB; - - int collidableIndexA = bodiesCPU[bodyIndexA].m_collidableIdx; - int collidableIndexB = bodiesCPU[bodyIndexB].m_collidableIdx; - - if (collidablesCPU[collidableIndexA].m_shapeType==SHAPE_CONCAVE_TRIMESH) + int numSubTrees = subTreesGPU->size(); + btVector3 bvhAabbMin = bvhData[0]->m_bvhAabbMin; + btVector3 bvhAabbMax = bvhData[0]->m_bvhAabbMax; + btVector3 bvhQuantization = bvhData[0]->m_bvhQuantization; { - //check aabbWS for bodyB against optimized BVH - btVector3 aabbMin = (const btVector3&)aabbsWSCPU[bodyIndexB].m_min[0]; - aabbMin[3] = 0.f; - btVector3 aabbMax = (const btVector3&)aabbsWSCPU[bodyIndexB].m_max[0]; - aabbMax[3] = 0.f; - bvhData[0]->reportAabbOverlappingNodex(&triCallback, aabbMin,aabbMax); + BT_PROFILE("m_bvhTraversalKernel"); + numConcavePairs = numConcavePairsOut.at(0); + btLauncherCL launcher(m_queue, m_bvhTraversalKernel); + launcher.setBuffer( pairs->getBufferCL()); + launcher.setBuffer( bodyBuf->getBufferCL()); + launcher.setBuffer( gpuCollidables.getBufferCL()); + launcher.setBuffer( clAabbsWS.getBufferCL()); + launcher.setBuffer( triangleConvexPairsOut.getBufferCL()); + launcher.setBuffer( numConcavePairsOut.getBufferCL()); + launcher.setBuffer( subTreesGPU->getBufferCL()); + launcher.setBuffer( treeNodesGPU->getBufferCL()); + launcher.setConst( bvhAabbMin); + launcher.setConst( bvhAabbMax); + launcher.setConst( bvhQuantization); + launcher.setConst(numSubTrees); + launcher.setConst( nPairs ); + launcher.setConst( maxTriConvexPairCapacity); + int num = nPairs; + launcher.launch1D( num); + clFinish(m_queue); + numConcavePairs = numConcavePairsOut.at(0); + + if (numConcavePairs > maxTriConvexPairCapacity) + { + static int exceeded_maxTriConvexPairCapacity_count = 0; + printf("Rxceeded %d times the maxTriConvexPairCapacity (found %d but max is %d)\n", exceeded_maxTriConvexPairCapacity_count++, + numConcavePairs,maxTriConvexPairCapacity); + numConcavePairs = maxTriConvexPairCapacity; + } + triangleConvexPairsOut.resize(numConcavePairs); + if (numConcavePairs) + { + //now perform a SAT test for each triangle-convex element (stored in triangleConvexPairsOut) + BT_PROFILE("findConcaveSeparatingAxisKernel"); + btBufferInfoCL bInfo[] = { + btBufferInfoCL( triangleConvexPairsOut.getBufferCL() ), + btBufferInfoCL( bodyBuf->getBufferCL(),true), + btBufferInfoCL( gpuCollidables.getBufferCL(),true), + btBufferInfoCL( convexData.getBufferCL(),true), + btBufferInfoCL( gpuVertices.getBufferCL(),true), + btBufferInfoCL( gpuUniqueEdges.getBufferCL(),true), + btBufferInfoCL( gpuFaces.getBufferCL(),true), + btBufferInfoCL( gpuIndices.getBufferCL(),true), + btBufferInfoCL( clAabbsWS.getBufferCL(),true), + btBufferInfoCL( sepNormals.getBufferCL()), + btBufferInfoCL( hasSeparatingNormals.getBufferCL()), + btBufferInfoCL( concaveSepNormals.getBufferCL()) + }; + + btLauncherCL launcher(m_queue, m_findConcaveSeparatingAxisKernel); + launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(btBufferInfoCL) ); + + launcher.setConst( numConcavePairs ); + + int num = numConcavePairs; + launcher.launch1D( num); + clFinish(m_queue); + } } } - - //now perform the tree query on GPU - - int numNodes = bvhData[0]->getLeafNodeArray().size(); - btOpenCLArray treeNodesGPU(this->m_context,this->m_queue,numNodes); - treeNodesGPU.copyFromHost(bvhData[0]->getQuantizedNodeArray()); - int numSubTrees = bvhData[0]->getSubtreeInfoArray().size(); - btOpenCLArray subTreesGPU(this->m_context,this->m_queue,numSubTrees); - subTreesGPU.copyFromHost(bvhData[0]->getSubtreeInfoArray()); - - - - btVector3 bvhAabbMin = bvhData[0]->m_bvhAabbMin; - btVector3 bvhAabbMax = bvhData[0]->m_bvhAabbMax; - btVector3 bvhQuantization = bvhData[0]->m_bvhQuantization; - { - int np = numConcavePairsOut.at(0); - printf("np=%d\n", np); - btLauncherCL launcher(m_queue, m_bvhTraversalKernel); - launcher.setBuffer( pairs->getBufferCL()); - launcher.setBuffer( bodyBuf->getBufferCL()); - launcher.setBuffer( gpuCollidables.getBufferCL()); - launcher.setBuffer( clAabbsWS.getBufferCL()); - launcher.setBuffer( triangleConvexPairsOut.getBufferCL()); - launcher.setBuffer( numConcavePairsOut.getBufferCL()); - launcher.setBuffer( subTreesGPU.getBufferCL()); - launcher.setBuffer( treeNodesGPU.getBufferCL()); - launcher.setConst( bvhAabbMin); - launcher.setConst( bvhAabbMax); - launcher.setConst( bvhQuantization); - launcher.setConst(numSubTrees); - launcher.setConst( nPairs ); - launcher.setConst( maxTriConvexPairCapacity); - int num = nPairs; - launcher.launch1D( num); - clFinish(m_queue); - np = numConcavePairsOut.at(0); - triangleConvexPairsOut.resize(np); - btAlignedObjectArray pairsOutCPU; - triangleConvexPairsOut.copyToHost(pairsOutCPU); - clFinish(m_queue); - - printf("np=%d\n", np); - - } - printf("-----------------------\n", numConcave); } - - { BT_PROFILE("findCompoundPairsKernel"); btBufferInfoCL bInfo[] = @@ -467,7 +440,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray& worldVertsA1GPU, btOpenCLArray& worldVertsB2GPU, btAlignedObjectArray& bvhData, + btOpenCLArray* treeNodesGPU, + btOpenCLArray* subTreesGPU, int numObjects, int maxTriConvexPairCapacity, btOpenCLArray& triangleConvexPairs, diff --git a/opencl/gpu_sat/kernels/sat.cl b/opencl/gpu_sat/kernels/sat.cl index 431834421..3417daf91 100644 --- a/opencl/gpu_sat/kernels/sat.cl +++ b/opencl/gpu_sat/kernels/sat.cl @@ -939,11 +939,7 @@ __kernel void findSeparatingAxisKernel( __global const int2* pairs, __global btAabbCL* aabbs, __global volatile float4* separatingNormals, __global volatile int* hasSeparatingAxis, - __global int4* concavePairsOut, - __global float4* concaveSeparatingNormalsOut, - __global volatile int* numConcavePairsOut, - int numPairs, - int maxNumConcavePairsCapacity + int numPairs ) { @@ -970,230 +966,6 @@ __kernel void findSeparatingAxisKernel( __global const int2* pairs, return; } - if ((collidables[collidableIndexA].m_shapeType==SHAPE_CONCAVE_TRIMESH))// && (collidables[collidableIndexB].m_shapeType==SHAPE_CONVEX_HULL)) - { - - int numFacesA = convexShapes[shapeIndexA].m_numFaces; - int numActualConcaveConvexTests = 0; - - for (int f=0;f vert.x) - triAabb.m_min.x = vert.x; - if (triAabb.m_min.y > vert.y) - triAabb.m_min.y = vert.y; - if (triAabb.m_min.z > vert.z) - triAabb.m_min.z = vert.z; - - if (triAabb.m_max.x < vert.x) - triAabb.m_max.x = vert.x; - if (triAabb.m_max.y < vert.y) - triAabb.m_max.y = vert.y; - if (triAabb.m_max.z < vert.z) - triAabb.m_max.z = vert.z; -#else - triAabb.m_min = min(triAabb.m_min,vert); - triAabb.m_max = max(triAabb.m_max,vert); -#endif - } - - overlap = true; - overlap = (triAabb.m_min.x > aabbs[bodyIndexB].m_max.x || triAabb.m_max.x < aabbs[bodyIndexB].m_min.x) ? false : overlap; - overlap = (triAabb.m_min.z > aabbs[bodyIndexB].m_max.z || triAabb.m_max.z < aabbs[bodyIndexB].m_min.z) ? false : overlap; - overlap = (triAabb.m_min.y > aabbs[bodyIndexB].m_max.y || triAabb.m_max.y < aabbs[bodyIndexB].m_min.y) ? false : overlap; - - if (overlap) - { - float dmin = FLT_MAX; - int hasSeparatingAxis=5; - float4 sepAxis=make_float4(1,2,3,4); - -#if 1 - - int localCC=0; - numActualConcaveConvexTests++; - - //a triangle has 3 unique edges - convexPolyhedronA.m_numUniqueEdges = 3; - convexPolyhedronA.m_uniqueEdgesOffset = 0; - float4 uniqueEdgesA[3]; - - uniqueEdgesA[0] = (verticesA[1]-verticesA[0]); - uniqueEdgesA[1] = (verticesA[2]-verticesA[1]); - uniqueEdgesA[2] = (verticesA[0]-verticesA[2]); - - - convexPolyhedronA.m_faceOffset = 0; - - float4 normal = make_float4(face.m_plane.x,face.m_plane.y,face.m_plane.z,0.f); - - btGpuFace facesA[TRIANGLE_NUM_CONVEX_FACES]; - int indicesA[3+3+2+2+2]; - int curUsedIndices=0; - int fidx=0; - - //front size of triangle - { - facesA[fidx].m_indexOffset=curUsedIndices; - indicesA[0] = 0; - indicesA[1] = 1; - indicesA[2] = 2; - curUsedIndices+=3; - float c = face.m_plane.w; - facesA[fidx].m_plane.x = normal.x; - facesA[fidx].m_plane.y = normal.y; - facesA[fidx].m_plane.z = normal.z; - facesA[fidx].m_plane.w = c; - facesA[fidx].m_numIndices=3; - } - fidx++; - //back size of triangle - { - facesA[fidx].m_indexOffset=curUsedIndices; - indicesA[3]=2; - indicesA[4]=1; - indicesA[5]=0; - curUsedIndices+=3; - float c = dot(normal,verticesA[0]); - float c1 = -face.m_plane.w; - facesA[fidx].m_plane.x = -normal.x; - facesA[fidx].m_plane.y = -normal.y; - facesA[fidx].m_plane.z = -normal.z; - facesA[fidx].m_plane.w = c; - facesA[fidx].m_numIndices=3; - } - fidx++; - - bool addEdgePlanes = true; - if (addEdgePlanes) - { - int numVertices=3; - int prevVertex = numVertices-1; - for (int i=0;i=numConcavePairs) + return; + int pairIdx = i; + + int bodyIndexA = concavePairs[i].x; + int bodyIndexB = concavePairs[i].y; + + int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx; + int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx; + + int shapeIndexA = collidables[collidableIndexA].m_shapeIndex; + int shapeIndexB = collidables[collidableIndexB].m_shapeIndex; + + + int numFacesA = convexShapes[shapeIndexA].m_numFaces; + int numActualConcaveConvexTests = 0; + + int f = concavePairs[i].z; + + bool overlap = false; + + ConvexPolyhedronCL convexPolyhedronA; + + //add 3 vertices of the triangle + convexPolyhedronA.m_numVertices = 3; + convexPolyhedronA.m_vertexOffset = 0; + float4 localCenter = make_float4(0.f,0.f,0.f,0.f); + + btGpuFace face = faces[convexShapes[shapeIndexA].m_faceOffset+f]; + float4 triMinAabb, triMaxAabb; + btAabbCL triAabb; + triAabb.m_min = make_float4(1e30f,1e30f,1e30f,0.f); + triAabb.m_max = make_float4(-1e30f,-1e30f,-1e30f,0.f); + + float4 verticesA[3]; + for (int i=0;i<3;i++) + { + int index = indices[face.m_indexOffset+i]; + float4 vert = vertices[convexShapes[shapeIndexA].m_vertexOffset+index]; + verticesA[i] = vert; + localCenter += vert; + + triAabb.m_min = min(triAabb.m_min,vert); + triAabb.m_max = max(triAabb.m_max,vert); + + } + + overlap = true; + overlap = (triAabb.m_min.x > aabbs[bodyIndexB].m_max.x || triAabb.m_max.x < aabbs[bodyIndexB].m_min.x) ? false : overlap; + overlap = (triAabb.m_min.z > aabbs[bodyIndexB].m_max.z || triAabb.m_max.z < aabbs[bodyIndexB].m_min.z) ? false : overlap; + overlap = (triAabb.m_min.y > aabbs[bodyIndexB].m_max.y || triAabb.m_max.y < aabbs[bodyIndexB].m_min.y) ? false : overlap; + + if (overlap) + { + float dmin = FLT_MAX; + int hasSeparatingAxis=5; + float4 sepAxis=make_float4(1,2,3,4); + + int localCC=0; + numActualConcaveConvexTests++; + + //a triangle has 3 unique edges + convexPolyhedronA.m_numUniqueEdges = 3; + convexPolyhedronA.m_uniqueEdgesOffset = 0; + float4 uniqueEdgesA[3]; + + uniqueEdgesA[0] = (verticesA[1]-verticesA[0]); + uniqueEdgesA[1] = (verticesA[2]-verticesA[1]); + uniqueEdgesA[2] = (verticesA[0]-verticesA[2]); + + + convexPolyhedronA.m_faceOffset = 0; + + float4 normal = make_float4(face.m_plane.x,face.m_plane.y,face.m_plane.z,0.f); + + btGpuFace facesA[TRIANGLE_NUM_CONVEX_FACES]; + int indicesA[3+3+2+2+2]; + int curUsedIndices=0; + int fidx=0; + + //front size of triangle + { + facesA[fidx].m_indexOffset=curUsedIndices; + indicesA[0] = 0; + indicesA[1] = 1; + indicesA[2] = 2; + curUsedIndices+=3; + float c = face.m_plane.w; + facesA[fidx].m_plane.x = normal.x; + facesA[fidx].m_plane.y = normal.y; + facesA[fidx].m_plane.z = normal.z; + facesA[fidx].m_plane.w = c; + facesA[fidx].m_numIndices=3; + } + fidx++; + //back size of triangle + { + facesA[fidx].m_indexOffset=curUsedIndices; + indicesA[3]=2; + indicesA[4]=1; + indicesA[5]=0; + curUsedIndices+=3; + float c = dot(normal,verticesA[0]); + float c1 = -face.m_plane.w; + facesA[fidx].m_plane.x = -normal.x; + facesA[fidx].m_plane.y = -normal.y; + facesA[fidx].m_plane.z = -normal.z; + facesA[fidx].m_plane.w = c; + facesA[fidx].m_numIndices=3; + } + fidx++; + + bool addEdgePlanes = true; + if (addEdgePlanes) + { + int numVertices=3; + int prevVertex = numVertices-1; + for (int i=0;i vert.x)\n" -" triAabb.m_min.x = vert.x;\n" -" if (triAabb.m_min.y > vert.y)\n" -" triAabb.m_min.y = vert.y;\n" -" if (triAabb.m_min.z > vert.z)\n" -" triAabb.m_min.z = vert.z;\n" -"\n" -" if (triAabb.m_max.x < vert.x)\n" -" triAabb.m_max.x = vert.x;\n" -" if (triAabb.m_max.y < vert.y)\n" -" triAabb.m_max.y = vert.y;\n" -" if (triAabb.m_max.z < vert.z)\n" -" triAabb.m_max.z = vert.z;\n" -"#else \n" -" triAabb.m_min = min(triAabb.m_min,vert); \n" -" triAabb.m_max = max(triAabb.m_max,vert); \n" -"#endif \n" -" }\n" -"\n" -" overlap = true;\n" -" overlap = (triAabb.m_min.x > aabbs[bodyIndexB].m_max.x || triAabb.m_max.x < aabbs[bodyIndexB].m_min.x) ? false : overlap;\n" -" overlap = (triAabb.m_min.z > aabbs[bodyIndexB].m_max.z || triAabb.m_max.z < aabbs[bodyIndexB].m_min.z) ? false : overlap;\n" -" overlap = (triAabb.m_min.y > aabbs[bodyIndexB].m_max.y || triAabb.m_max.y < aabbs[bodyIndexB].m_min.y) ? false : overlap;\n" -" \n" -" if (overlap)\n" -" {\n" -" float dmin = FLT_MAX;\n" -" int hasSeparatingAxis=5;\n" -" float4 sepAxis=make_float4(1,2,3,4);\n" -"\n" -"#if 1\n" -" \n" -" int localCC=0;\n" -" numActualConcaveConvexTests++;\n" -"\n" -" //a triangle has 3 unique edges\n" -" convexPolyhedronA.m_numUniqueEdges = 3;\n" -" convexPolyhedronA.m_uniqueEdgesOffset = 0;\n" -" float4 uniqueEdgesA[3];\n" -" \n" -" uniqueEdgesA[0] = (verticesA[1]-verticesA[0]);\n" -" uniqueEdgesA[1] = (verticesA[2]-verticesA[1]);\n" -" uniqueEdgesA[2] = (verticesA[0]-verticesA[2]);\n" -"\n" -"\n" -" convexPolyhedronA.m_faceOffset = 0;\n" -" \n" -" float4 normal = make_float4(face.m_plane.x,face.m_plane.y,face.m_plane.z,0.f);\n" -" \n" -" btGpuFace facesA[TRIANGLE_NUM_CONVEX_FACES];\n" -" int indicesA[3+3+2+2+2];\n" -" int curUsedIndices=0;\n" -" int fidx=0;\n" -"\n" -" //front size of triangle\n" -" {\n" -" facesA[fidx].m_indexOffset=curUsedIndices;\n" -" indicesA[0] = 0;\n" -" indicesA[1] = 1;\n" -" indicesA[2] = 2;\n" -" curUsedIndices+=3;\n" -" float c = face.m_plane.w;\n" -" facesA[fidx].m_plane.x = normal.x;\n" -" facesA[fidx].m_plane.y = normal.y;\n" -" facesA[fidx].m_plane.z = normal.z;\n" -" facesA[fidx].m_plane.w = c;\n" -" facesA[fidx].m_numIndices=3;\n" -" }\n" -" fidx++;\n" -" //back size of triangle\n" -" {\n" -" facesA[fidx].m_indexOffset=curUsedIndices;\n" -" indicesA[3]=2;\n" -" indicesA[4]=1;\n" -" indicesA[5]=0;\n" -" curUsedIndices+=3;\n" -" float c = dot(normal,verticesA[0]);\n" -" float c1 = -face.m_plane.w;\n" -" facesA[fidx].m_plane.x = -normal.x;\n" -" facesA[fidx].m_plane.y = -normal.y;\n" -" facesA[fidx].m_plane.z = -normal.z;\n" -" facesA[fidx].m_plane.w = c;\n" -" facesA[fidx].m_numIndices=3;\n" -" }\n" -" fidx++;\n" -"\n" -" bool addEdgePlanes = true;\n" -" if (addEdgePlanes)\n" -" {\n" -" int numVertices=3;\n" -" int prevVertex = numVertices-1;\n" -" for (int i=0;i=numConcavePairs)\n" +" return;\n" +" int pairIdx = i;\n" +"\n" +" int bodyIndexA = concavePairs[i].x;\n" +" int bodyIndexB = concavePairs[i].y;\n" +"\n" +" int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n" +" int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n" +"\n" +" int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;\n" +" int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;\n" +"\n" +"\n" +" int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n" +" int numActualConcaveConvexTests = 0;\n" +" \n" +" int f = concavePairs[i].z;\n" +" \n" +" bool overlap = false;\n" +" \n" +" ConvexPolyhedronCL convexPolyhedronA;\n" +"\n" +" //add 3 vertices of the triangle\n" +" convexPolyhedronA.m_numVertices = 3;\n" +" convexPolyhedronA.m_vertexOffset = 0;\n" +" float4 localCenter = make_float4(0.f,0.f,0.f,0.f);\n" +"\n" +" btGpuFace face = faces[convexShapes[shapeIndexA].m_faceOffset+f];\n" +" float4 triMinAabb, triMaxAabb;\n" +" btAabbCL triAabb;\n" +" triAabb.m_min = make_float4(1e30f,1e30f,1e30f,0.f);\n" +" triAabb.m_max = make_float4(-1e30f,-1e30f,-1e30f,0.f);\n" +" \n" +" float4 verticesA[3];\n" +" for (int i=0;i<3;i++)\n" +" {\n" +" int index = indices[face.m_indexOffset+i];\n" +" float4 vert = vertices[convexShapes[shapeIndexA].m_vertexOffset+index];\n" +" verticesA[i] = vert;\n" +" localCenter += vert;\n" +" \n" +" triAabb.m_min = min(triAabb.m_min,vert); \n" +" triAabb.m_max = max(triAabb.m_max,vert); \n" +"\n" +" }\n" +"\n" +" overlap = true;\n" +" overlap = (triAabb.m_min.x > aabbs[bodyIndexB].m_max.x || triAabb.m_max.x < aabbs[bodyIndexB].m_min.x) ? false : overlap;\n" +" overlap = (triAabb.m_min.z > aabbs[bodyIndexB].m_max.z || triAabb.m_max.z < aabbs[bodyIndexB].m_min.z) ? false : overlap;\n" +" overlap = (triAabb.m_min.y > aabbs[bodyIndexB].m_max.y || triAabb.m_max.y < aabbs[bodyIndexB].m_min.y) ? false : overlap;\n" +" \n" +" if (overlap)\n" +" {\n" +" float dmin = FLT_MAX;\n" +" int hasSeparatingAxis=5;\n" +" float4 sepAxis=make_float4(1,2,3,4);\n" +"\n" +" int localCC=0;\n" +" numActualConcaveConvexTests++;\n" +"\n" +" //a triangle has 3 unique edges\n" +" convexPolyhedronA.m_numUniqueEdges = 3;\n" +" convexPolyhedronA.m_uniqueEdgesOffset = 0;\n" +" float4 uniqueEdgesA[3];\n" +" \n" +" uniqueEdgesA[0] = (verticesA[1]-verticesA[0]);\n" +" uniqueEdgesA[1] = (verticesA[2]-verticesA[1]);\n" +" uniqueEdgesA[2] = (verticesA[0]-verticesA[2]);\n" +"\n" +"\n" +" convexPolyhedronA.m_faceOffset = 0;\n" +" \n" +" float4 normal = make_float4(face.m_plane.x,face.m_plane.y,face.m_plane.z,0.f);\n" +" \n" +" btGpuFace facesA[TRIANGLE_NUM_CONVEX_FACES];\n" +" int indicesA[3+3+2+2+2];\n" +" int curUsedIndices=0;\n" +" int fidx=0;\n" +"\n" +" //front size of triangle\n" +" {\n" +" facesA[fidx].m_indexOffset=curUsedIndices;\n" +" indicesA[0] = 0;\n" +" indicesA[1] = 1;\n" +" indicesA[2] = 2;\n" +" curUsedIndices+=3;\n" +" float c = face.m_plane.w;\n" +" facesA[fidx].m_plane.x = normal.x;\n" +" facesA[fidx].m_plane.y = normal.y;\n" +" facesA[fidx].m_plane.z = normal.z;\n" +" facesA[fidx].m_plane.w = c;\n" +" facesA[fidx].m_numIndices=3;\n" +" }\n" +" fidx++;\n" +" //back size of triangle\n" +" {\n" +" facesA[fidx].m_indexOffset=curUsedIndices;\n" +" indicesA[3]=2;\n" +" indicesA[4]=1;\n" +" indicesA[5]=0;\n" +" curUsedIndices+=3;\n" +" float c = dot(normal,verticesA[0]);\n" +" float c1 = -face.m_plane.w;\n" +" facesA[fidx].m_plane.x = -normal.x;\n" +" facesA[fidx].m_plane.y = -normal.y;\n" +" facesA[fidx].m_plane.z = -normal.z;\n" +" facesA[fidx].m_plane.w = c;\n" +" facesA[fidx].m_numIndices=3;\n" +" }\n" +" fidx++;\n" +"\n" +" bool addEdgePlanes = true;\n" +" if (addEdgePlanes)\n" +" {\n" +" int numVertices=3;\n" +" int prevVertex = numVertices-1;\n" +" for (int i=0;i