add support for BVH acceleration for concave trianglemesh collision against convex hulls

bugfix/improvement in batching

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)
 	{

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

demo/gpudemo/GpuDemo.h

@@ -34,13 +34,13 @@ public:
                     :useOpenCL(true),
                     preferredOpenCLPlatformIndex(-1),
                     preferredOpenCLDeviceIndex(-1),
-					arraySizeX(5),
+					arraySizeX(30),
-		arraySizeY(5 ),
+		arraySizeY(10 ),
-		arraySizeZ(5),
+		arraySizeZ(30),
 		m_useConcaveMesh(false),
-		gapX(4.3),
+		gapX(6.3),
-		gapY(2.0),
+		gapY(12.0),
-		gapZ(4.3),
+		gapZ(6.3),
                     m_instancingRenderer(0),
 					m_window(0)
             {

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();
 			}

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/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";
 
 	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<btVector3> verts;
 		for (int i=0;i<shape->m_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);
 
 }

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);

opencl/gpu_rigidbody/host/Solver.cpp

@@ -818,9 +818,10 @@ void Solver::sortContacts(  const btOpenCLArray<btRigidBodyCL>* bodyBuf,
 
 */
 
-void Solver::batchContacts(  btOpenCLArray<btContact4>* contacts, int nContacts, btOpenCLArray<unsigned int>* nNative, btOpenCLArray<unsigned int>* offsetsNative, int staticIdx )
+void	Solver::batchContacts(  btOpenCLArray<btContact4>* contacts, int nContacts, btOpenCLArray<unsigned int>* nNative, btOpenCLArray<unsigned int>* offsetsNative, int staticIdx )
 {
-
+	
+	int numWorkItems = 64*N_SPLIT*N_SPLIT;
 	{
 		BT_PROFILE("batch generation");
 		
@@ -829,7 +830,7 @@ void Solver::batchContacts(  btOpenCLArray<btContact4>* 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<SolverDebugInfo> gpuDebugInfo(data->m_device,numWorkItems);
@@ -837,12 +838,14 @@ void Solver::batchContacts(  btOpenCLArray<btContact4>* 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<btContact4>* contacts, int nContacts,
 	btAssert(m_contactBuffer->size()==nContacts);
 	//contacts->copyFromOpenCLArray( *m_contactBuffer);
 	//clFinish(m_queue);//needed?
-
+	
+	
 	
 }
 

opencl/gpu_rigidbody/host/Solver.h

@@ -132,7 +132,7 @@ class Solver : public SolverBase
 			btOpenCLArray<btContact4>* contactsIn, btOpenCLArray<btGpuConstraint4>* contactCOut, void* additionalData, 
 			int nContacts, const ConstraintCfg& cfg );
 
-		void batchContacts( btOpenCLArray<btContact4>* contacts, int nContacts, btOpenCLArray<unsigned int>* n, btOpenCLArray<unsigned int>* offsets, int staticIdx );
+		void	batchContacts( btOpenCLArray<btContact4>* contacts, int nContacts, btOpenCLArray<unsigned int>* n, btOpenCLArray<unsigned int>* offsets, int staticIdx );
 
 };
 

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;
 	}

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_averageExtent = .2f;//@TODO m_averageObjExtent;
-        csCfg.m_staticIdx = -1;//m_static0Index;//m_planeBodyIndex;
+        csCfg.m_staticIdx = 0;//m_static0Index;//m_planeBodyIndex;
         
         btOpenCLArray<btContact4>* contactsIn = m_data->m_pBufContactOutGPU;
         btOpenCLArray<btRigidBodyCL>* 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)
+					if (gpuBatchContacts)
-                {
+					{
-					maxNumBatches=250;//for now
+						BT_PROFILE("gpu batchContacts");
-                    BT_PROFILE("gpu batchContacts");
+						maxNumBatches = 50;
-                    m_data->m_solverGPU->batchContacts( (btOpenCLArray<btContact4>*)contactNative, nContacts, m_data->m_solverGPU->m_numConstraints, m_data->m_solverGPU->m_offsets, csCfg.m_staticIdx );
+						m_data->m_solverGPU->batchContacts( (btOpenCLArray<btContact4>*)contactNative, nContacts, m_data->m_solverGPU->m_numConstraints, m_data->m_solverGPU->m_offsets, csCfg.m_staticIdx );
-                } else
+					} else
-                {
+					{
-                    BT_PROFILE("cpu batchContacts");
+						BT_PROFILE("cpu batchContacts");
-                    btAlignedObjectArray<btContact4> cpuContacts;
+						btAlignedObjectArray<btContact4> cpuContacts;
-                    btOpenCLArray<btContact4>* contactsIn = m_data->m_pBufContactOutGPU;
+						btOpenCLArray<btContact4>* contactsIn = m_data->m_pBufContactOutGPU;
-                    contactsIn->copyToHost(cpuContacts);
+						contactsIn->copyToHost(cpuContacts);
                     
-                    btOpenCLArray<unsigned int>* countsNative = m_data->m_solverGPU->m_numConstraints;
+						btOpenCLArray<unsigned int>* countsNative = m_data->m_solverGPU->m_numConstraints;
-                    btOpenCLArray<unsigned int>* offsetsNative = m_data->m_solverGPU->m_offsets;
+						btOpenCLArray<unsigned int>* offsetsNative = m_data->m_solverGPU->m_offsets;
                     
-                    btAlignedObjectArray<unsigned int> nNativeHost;
+						btAlignedObjectArray<unsigned int> nNativeHost;
-                    btAlignedObjectArray<unsigned int> offsetsNativeHost;
+						btAlignedObjectArray<unsigned int> offsetsNativeHost;
                     
-                    {
+						{
-                        BT_PROFILE("countsNative/offsetsNative copyToHost");
+							BT_PROFILE("countsNative/offsetsNative copyToHost");
-                        countsNative->copyToHost(nNativeHost);
+							countsNative->copyToHost(nNativeHost);
-                        offsetsNative->copyToHost(offsetsNativeHost);
+							offsetsNative->copyToHost(offsetsNativeHost);
-                    }
+						}
                     
                     
-                    int numNonzeroGrid=0;
+						int numNonzeroGrid=0;
                     
-                    {
+						{
-                        BT_PROFILE("batch grid");
+							BT_PROFILE("batch grid");
-                        for(int i=0; i<BT_SOLVER_N_SPLIT*BT_SOLVER_N_SPLIT; i++)
+							for(int i=0; i<BT_SOLVER_N_SPLIT*BT_SOLVER_N_SPLIT; i++)
-                        {
+							{
-                            int n = (nNativeHost)[i];
+								int n = (nNativeHost)[i];
-                            int offset = (offsetsNativeHost)[i];
+								int offset = (offsetsNativeHost)[i];
                             
-                            if( n )
+								if( n )
-                            {
+								{
-                                numNonzeroGrid++;
+									numNonzeroGrid++;
-                                //printf("cpu batch\n");
+									//printf("cpu batch\n");
                                 
-                                int simdWidth = -1;
+									int simdWidth = -1;
-                                int numBatches = sortConstraintByBatch( &cpuContacts[0]+offset, n, simdWidth,csCfg.m_staticIdx );	//	on GPU
+									int numBatches = sortConstraintByBatch( &cpuContacts[0]+offset, n, simdWidth,csCfg.m_staticIdx ,numBodies);	//	on GPU
-                                maxNumBatches = btMax(numBatches,maxNumBatches);
+									maxNumBatches = btMax(numBatches,maxNumBatches);
                                 
-                                clFinish(m_data->m_queue);
+									clFinish(m_data->m_queue);
                                 
-                            }
+								}
-                        }
+							}
-                    }
+						}
-                    {
+						{
-                        BT_PROFILE("m_contactBuffer->copyFromHost");
+							BT_PROFILE("m_contactBuffer->copyFromHost");
-                        m_data->m_solverGPU->m_contactBuffer->copyFromHost((btAlignedObjectArray<btContact4>&)cpuContacts);
+							m_data->m_solverGPU->m_contactBuffer->copyFromHost((btAlignedObjectArray<btContact4>&)cpuContacts);
-                    }
+						}
-                  //  printf("maxNumBatches = %d\n", maxNumBatches);
+						
-				} 
+					} 
+					
                 }
                 
                 
-                
+                //printf("maxNumBatches = %d\n", maxNumBatches);
                 
                 if (nContacts)
                 {
@@ -704,8 +705,14 @@ btAlignedObjectArray<btSortData> sortData;
 btAlignedObjectArray<btContact4> 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<int> bodyUsed;
+	bodyUsed.resize(numBodies);
+	for (int q=0;q<numBodies;q++)
+		bodyUsed[q]=0;
+
+	BT_PROFILE("sortConstraintByBatch");
 	int numIter = 0;
     
 	sortData.resize(n);
@@ -747,12 +754,7 @@ inline int btGpuBatchingPgsSolver::sortConstraintByBatch( btContact4* cs, int n,
 				int bodyAS = cs[idx].m_bodyAPtrAndSignBit;
 				int bodyBS = cs[idx].m_bodyBPtrAndSignBit;
                 
-				/*if (bodyAS<0)
+				
-					printf("A static\n");
-                
-				if (bodyBS<0)
-					printf("B static\n");
-                */
                 
 				int bodyA = abs(bodyAS);
 				int bodyB = abs(bodyBS);
@@ -763,14 +765,20 @@ inline int btGpuBatchingPgsSolver::sortConstraintByBatch( btContact4* cs, int n,
 				unsigned int aUnavailable = flg[ aIdx/32 ] & (1<<(aIdx&31));
 				unsigned int bUnavailable = flg[ bIdx/32 ] & (1<<(bIdx&31));
                 
+				bool aIsStatic = (bodyAS<0) || bodyAS==staticIdx;
+				bool bIsStatic = (bodyBS<0) || bodyBS==staticIdx;
+
                 //use inv_mass!
-				aUnavailable = (bodyAS>=0)&&bodyAS!=staticIdx? aUnavailable:0;//
+				aUnavailable = !aIsStatic? aUnavailable:0;//
-				bUnavailable = (bodyBS>=0)&&bodyBS!=staticIdx? bUnavailable:0;
+				bUnavailable = !bIsStatic? bUnavailable:0;
                 
 				if( aUnavailable==0 && bUnavailable==0 ) // ok
 				{
-					flg[ aIdx/32 ] |= (1<<(aIdx&31));
+					if (!!aIsStatic)
-					flg[ bIdx/32 ] |= (1<<(bIdx&31));
+						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;

opencl/gpu_rigidbody/host/btGpuBatchingPgsSolver.h

@@ -15,7 +15,7 @@ protected:
 	struct btGpuBatchingPgsSolverInternalData*		m_data;
 
 	void batchContacts( btOpenCLArray<btContact4>* contacts, int nContacts, btOpenCLArray<unsigned int>* n, btOpenCLArray<unsigned int>* 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<btRigidBodyCL>* bodyBuf, const btOpenCLArray<btInertiaCL>* shapeBuf, 
 			btOpenCLArray<btGpuConstraint4>* constraint, void* additionalData, int n ,int maxNumBatches, int numIterations);
 

opencl/gpu_rigidbody/host/btGpuNarrowPhase.cpp

@@ -63,6 +63,10 @@ struct btGpuNarrowPhaseInternalData
 	btAlignedObjectArray<btSapAabb>* m_localShapeAABBCPU;
 
 	btAlignedObjectArray<class btOptimizedBvh*> m_bvhData;
+	btOpenCLArray<btQuantizedBvhNode>*	m_treeNodesGPU;
+	btOpenCLArray<btBvhSubtreeInfo>*	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<Constraint4>(ctx,queue,config.m_maxBroadphasePairs,false);
 	//m_data->m_frictionCGPU = new btOpenCLArray<adl::Solver<adl::TYPE_CL>::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<btVector3>* vertices, btAlignedObjectArray<int>* 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<btVector3>* 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<btQuantizedBvhNode>*	treeNodesGPU = new btOpenCLArray<btQuantizedBvhNode>(this->m_context,this->m_queue,numNodes);
+	treeNodesGPU->copyFromHost(bvh->getQuantizedNodeArray());
+
+	int numSubTrees = bvh->getSubtreeInfoArray().size();
+	btOpenCLArray<btBvhSubtreeInfo>* subTreesGPU = new btOpenCLArray<btBvhSubtreeInfo>(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<btVector3>*
 	{
 		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<btInt4> 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,

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;
+							bool aStatic = (e.m_a<0) ||(ea==m_staticIdx);
-							bAvailable = (e.m_b<0)? 1: bAvailable;
+							bool bStatic = (e.m_b<0) ||(eb==m_staticIdx);
 							
-							aAvailable = (e.m_a==m_staticIdx)? 1: aAvailable;
+							aAvailable = aStatic? 1: aAvailable;
-							bAvailable = (e.m_b==m_staticIdx)? 1: bAvailable;
+							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;
 						}

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"
+"							bool aStatic = (e.m_a<0) ||(ea==m_staticIdx);\n"
-"							bAvailable = (e.m_b<0)? 1: bAvailable;\n"
+"							bool bStatic = (e.m_b<0) ||(eb==m_staticIdx);\n"
 "							\n"
-"							aAvailable = (e.m_a==m_staticIdx)? 1: aAvailable;\n"
+"							aAvailable = aStatic? 1: aAvailable;\n"
-"							bAvailable = (e.m_b==m_staticIdx)? 1: bAvailable;\n"
+"							bAvailable = bStatic? 1: bAvailable;\n"
 "\n"
 "							bool success = (aAvailable && bAvailable);\n"
 "							if(success)\n"
 "							{\n"
-"								writeBuf( ldsFixedBuffer, ea );\n"
+"							\n"
-"								writeBuf( ldsFixedBuffer, eb );\n"
+"								if (!aStatic)\n"
+"									writeBuf( ldsFixedBuffer, ea );\n"
+"								if (!bStatic)\n"
+"									writeBuf( ldsFixedBuffer, eb );\n"
 "							}\n"
 "							done = success;\n"
 "						}\n"

opencl/gpu_sat/host/ConvexHullContact.cpp

@@ -23,7 +23,7 @@ subject to the following restrictions:
 #include "ConvexHullContact.h"
 #include <string.h>//memcpy
 #include "btConvexPolyhedronCL.h"
-#include "btOptimizedBvh.h"
+
 
 typedef btAlignedObjectArray<btVector3> 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<btI
             btOpenCLArray<btVector3>& worldVertsA1GPU,
             btOpenCLArray<btVector3>& worldVertsB2GPU,    
 			btAlignedObjectArray<class btOptimizedBvh*>& bvhData,
+			btOpenCLArray<btQuantizedBvhNode>*	treeNodesGPU,
+			btOpenCLArray<btBvhSubtreeInfo>*	subTreesGPU,
 			int numObjects,
 			int maxTriConvexPairCapacity,
 			btOpenCLArray<btInt4>& triangleConvexPairsOut,
@@ -231,146 +238,112 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 	int numCompoundPairs = 0;
 
 	bool findSeparatingAxisOnGpu = true;//false;
-	int numConcave =0;
+	int numConcavePairs =0;
 
 	{
 		clFinish(m_queue);
 		if (findSeparatingAxisOnGpu)
 		{
-	
 		
-			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()),
-				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)
+				BT_PROFILE("findSeparatingAxisKernel");
-					numConcave = maxTriConvexPairCapacity;
+				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);
+				btLauncherCL launcher(m_queue, m_findSeparatingAxisKernel);
-				btAlignedObjectArray<btInt4> triangleConvexPairsOutCPU;
+				launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(btBufferInfoCL) );
-				triangleConvexPairsOut.copyToHost(triangleConvexPairsOutCPU);
+				launcher.setConst( nPairs  );
-				printf("-----------------------\n", numConcave);
-				printf("got %d concave pairs\n", numConcave);
-				btAssert(numConcave = triangleConvexPairsOutCPU.size());
 
-				for (int i=0;i<triangleConvexPairsOutCPU.size();i++)
+				int num = nPairs;
-				{
+				launcher.launch1D( num);
-					printf("bodyIndexA = %d, bodyIndexB = %d\n", triangleConvexPairsOutCPU[i].x,triangleConvexPairsOutCPU[i].y);
+				clFinish(m_queue);
-					printf("triangleIndex = %d\n", triangleConvexPairsOutCPU[i].z);
+			}
-				}
-				printf("-----------------------\n", numConcave);
-				printf("Now using BVH query\n" );
-				btAlignedObjectArray<btCollidable> collidablesCPU;
-				gpuCollidables.copyToHost(collidablesCPU);
-				btAlignedObjectArray<btRigidBodyCL> bodiesCPU;
-				bodyBuf->copyToHost(bodiesCPU);
-				btAlignedObjectArray<btInt2> pairsCPU;
 
-				btAlignedObjectArray<btYetAnotherAabb> aabbsWSCPU;
+			//now perform the tree query on GPU
-				clAabbsWS.copyToHost(aabbsWSCPU);
+			{
-
-				pairs->copyToHost(pairsCPU);
-				MyTriangleCallback triCallback;
 				
-
+				int numNodes = bvhData.size()? bvhData[0]->getQuantizedNodeArray().size() : 0;
-				for (int i=0;i<pairsCPU.size();i++)
+				if (numNodes)
 				{
-					int bodyIndexA = pairsCPU[i].x;
+					int numSubTrees = subTreesGPU->size();
-					int bodyIndexB = pairsCPU[i].y;
+					btVector3 bvhAabbMin = bvhData[0]->m_bvhAabbMin;
-
+					btVector3 bvhAabbMax = bvhData[0]->m_bvhAabbMax;
-					triCallback.m_bodyIndexA = bodyIndexA;
+					btVector3 bvhQuantization = bvhData[0]->m_bvhQuantization;
-					triCallback.m_bodyIndexB = bodyIndexB;
-
-					int collidableIndexA = bodiesCPU[bodyIndexA].m_collidableIdx;
-					int collidableIndexB = bodiesCPU[bodyIndexB].m_collidableIdx;
-
-					if (collidablesCPU[collidableIndexA].m_shapeType==SHAPE_CONCAVE_TRIMESH)
 					{
-						//check aabbWS for bodyB against optimized BVH
+						BT_PROFILE("m_bvhTraversalKernel");
-						btVector3 aabbMin = (const btVector3&)aabbsWSCPU[bodyIndexB].m_min[0];
+						numConcavePairs = numConcavePairsOut.at(0);
-						aabbMin[3] = 0.f;
+						btLauncherCL launcher(m_queue, m_bvhTraversalKernel);
-						btVector3 aabbMax = (const btVector3&)aabbsWSCPU[bodyIndexB].m_max[0];
+						launcher.setBuffer( pairs->getBufferCL());
-						aabbMax[3] = 0.f;
+						launcher.setBuffer(  bodyBuf->getBufferCL());
-						bvhData[0]->reportAabbOverlappingNodex(&triCallback, aabbMin,aabbMax);
+						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<btQuantizedBvhNode>	treeNodesGPU(this->m_context,this->m_queue,numNodes);
-				treeNodesGPU.copyFromHost(bvhData[0]->getQuantizedNodeArray());
-				int numSubTrees = bvhData[0]->getSubtreeInfoArray().size();
-				btOpenCLArray<btBvhSubtreeInfo>	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<btInt4> 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<btI
 
 		//concave-convex contact clipping
 
-		if (numConcave)
+		if (numConcavePairs)
 		{
 			BT_PROFILE("clipHullHullConcaveConvexKernel");
 			nContacts = m_totalContactsOut.at(0);
@@ -486,8 +459,8 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 			};
 			btLauncherCL launcher(m_queue, m_clipHullHullConcaveConvexKernel);
 			launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(btBufferInfoCL) );
-			launcher.setConst( numConcave  );
+			launcher.setConst( numConcavePairs  );
-			int num = numConcave;
+			int num = numConcavePairs;
 			launcher.launch1D( num);
 			clFinish(m_queue);
 			nContacts = m_totalContactsOut.at(0);

opencl/gpu_sat/host/ConvexHullContact.h

@@ -11,6 +11,7 @@
 #include "btContact4.h"
 #include "parallel_primitives/host/btInt2.h"
 #include "parallel_primitives/host/btInt4.h"
+#include "btOptimizedBvh.h"
 
 //#include "../../dynamics/basic_demo/Stubs/ChNarrowPhase.h"
 
@@ -36,6 +37,7 @@ struct GpuSatCollision
 	cl_device_id			m_device;
 	cl_command_queue		m_queue;
 	cl_kernel				m_findSeparatingAxisKernel;
+	cl_kernel				m_findConcaveSeparatingAxisKernel;
 	cl_kernel				m_findCompoundPairsKernel;
 	cl_kernel				m_processCompoundPairsKernel;
 
@@ -76,6 +78,8 @@ struct GpuSatCollision
            btOpenCLArray<btVector3>& worldVertsA1GPU,
            btOpenCLArray<btVector3>& worldVertsB2GPU,
 		   btAlignedObjectArray<class btOptimizedBvh*>& bvhData,
+		   btOpenCLArray<btQuantizedBvhNode>*	treeNodesGPU,
+			btOpenCLArray<btBvhSubtreeInfo>*	subTreesGPU,
 			int numObjects,
 			int maxTriConvexPairCapacity,
 			btOpenCLArray<btInt4>& triangleConvexPairs,

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,
+																					int numPairs
-																					__global float4* concaveSeparatingNormalsOut,
-																					__global volatile int* numConcavePairsOut,
-																					int numPairs,
-																					int maxNumConcavePairsCapacity
 																					)
 {
 
@@ -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<numFacesA;f++)
-			{
-			
-				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;
-#if 0
-//just in case some implementation doesn't support component-wise min and max for float4
-					if (triAabb.m_min.x > 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<numVertices;i++)
-						{
-							float4 v0 = verticesA[i];
-							float4 v1 = verticesA[prevVertex];
-	                                                
-							float4 edgeNormal = normalize(cross(normal,v1-v0));
-							float c = -dot(edgeNormal,v0);
-	
-							facesA[fidx].m_numIndices = 2;
-							facesA[fidx].m_indexOffset=curUsedIndices;
-							indicesA[curUsedIndices++]=i;
-							indicesA[curUsedIndices++]=prevVertex;
-	                                                
-							facesA[fidx].m_plane.x = edgeNormal.x;
-							facesA[fidx].m_plane.y = edgeNormal.y;
-							facesA[fidx].m_plane.z = edgeNormal.z;
-							facesA[fidx].m_plane.w = c;
-							fidx++;
-							prevVertex = i;
-						}
-					}
-					convexPolyhedronA.m_numFaces = TRIANGLE_NUM_CONVEX_FACES;
-					convexPolyhedronA.m_localCenter = localCenter*(1.f/3.f);
-			
-
-					float4 posA = rigidBodies[bodyIndexA].m_pos;
-					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;
-					
-					bool sepA = findSeparatingAxisLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,
-																											rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,
-																											DeltaC2,
-																											verticesA,uniqueEdgesA,facesA,indicesA,
-																											vertices,uniqueEdges,faces,indices,
-																											&sepAxis,&dmin);
-					hasSeparatingAxis = 4;
-					if (!sepA)
-					{
-						hasSeparatingAxis = 0;
-					} else
-					{
-						bool sepB = findSeparatingAxisLocalB(	&convexShapes[shapeIndexB],&convexPolyhedronA,rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,
-																												rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,
-																												DeltaC2,
-																												vertices,uniqueEdges,faces,indices,
-																												verticesA,uniqueEdgesA,facesA,indicesA,
-																												&sepAxis,&dmin);
-			
-						if (!sepB)
-						{
-							hasSeparatingAxis = 0;
-						} else
-						{
-							bool sepEE = findSeparatingAxisEdgeEdgeLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,
-																												rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,
-																												DeltaC2,
-																												verticesA,uniqueEdgesA,facesA,indicesA,
-																												vertices,uniqueEdges,faces,indices,
-																												&sepAxis,&dmin);
-				
-							if (!sepEE)
-							{
-								hasSeparatingAxis = 0;
-							} else
-							{
-								hasSeparatingAxis = 1;
-							}
-						}
-					}	
-#endif
-					
-					if (hasSeparatingAxis)
-					{
-						int pairIdx = atomic_inc(numConcavePairsOut);
-						if (pairIdx<maxNumConcavePairsCapacity)
-						{
-							concavePairsOut[pairIdx].x = bodyIndexA;
-							concavePairsOut[pairIdx].y = bodyIndexB;
-							concavePairsOut[pairIdx].z = f;
-							concavePairsOut[pairIdx].w = 3;
-							sepAxis.w = dmin;
-							concaveSeparatingNormalsOut[pairIdx]=sepAxis;
-						}
-					}
-				}
-			}
-			//todo//??
-			hasSeparatingAxis[i] = 0;
-			return;
-		}		
-
-	
 
 		if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONVEX_HULL) ||(collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL))
 		{
@@ -1201,10 +973,12 @@ __kernel void   findSeparatingAxisKernel( __global const int2* pairs,
 			return;
 		}
 			
+		if ((collidables[collidableIndexA].m_shapeType==SHAPE_CONCAVE_TRIMESH))
+		{
+			hasSeparatingAxis[i] = 0;
+			return;
+		}
 
-
-
-		
 		int numFacesA = convexShapes[shapeIndexA].m_numFaces;
 
 		float dmin = FLT_MAX;
@@ -1262,4 +1036,237 @@ __kernel void   findSeparatingAxisKernel( __global const int2* pairs,
 		
 	}
 
-}
+}
+
+
+
+
+// work-in-progress
+__kernel void   findConcaveSeparatingAxisKernel( __global int4* concavePairs,
+																					__global const BodyData* rigidBodies,
+																					__global const btCollidableGpu* collidables,
+																					__global const ConvexPolyhedronCL* convexShapes, 
+																					__global const float4* vertices,
+																					__global const float4* uniqueEdges,
+																					__global const btGpuFace* faces,
+																					__global const int* indices,
+																					__global btAabbCL* aabbs,
+																					__global volatile float4* separatingNormals,
+																					__global volatile int* hasSeparatingAxis,
+																					__global float4* concaveSeparatingNormalsOut,
+																					int numConcavePairs
+																					)
+{
+
+	int i = get_global_id(0);
+	if (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<numVertices;i++)
+			{
+				float4 v0 = verticesA[i];
+				float4 v1 = verticesA[prevVertex];
+                                            
+				float4 edgeNormal = normalize(cross(normal,v1-v0));
+				float c = -dot(edgeNormal,v0);
+
+				facesA[fidx].m_numIndices = 2;
+				facesA[fidx].m_indexOffset=curUsedIndices;
+				indicesA[curUsedIndices++]=i;
+				indicesA[curUsedIndices++]=prevVertex;
+                                            
+				facesA[fidx].m_plane.x = edgeNormal.x;
+				facesA[fidx].m_plane.y = edgeNormal.y;
+				facesA[fidx].m_plane.z = edgeNormal.z;
+				facesA[fidx].m_plane.w = c;
+				fidx++;
+				prevVertex = i;
+			}
+		}
+		convexPolyhedronA.m_numFaces = TRIANGLE_NUM_CONVEX_FACES;
+		convexPolyhedronA.m_localCenter = localCenter*(1.f/3.f);
+
+
+		float4 posA = rigidBodies[bodyIndexA].m_pos;
+		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;
+		
+		bool sepA = findSeparatingAxisLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,
+																								rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,
+																								DeltaC2,
+																								verticesA,uniqueEdgesA,facesA,indicesA,
+																								vertices,uniqueEdges,faces,indices,
+																								&sepAxis,&dmin);
+		hasSeparatingAxis = 4;
+		if (!sepA)
+		{
+			hasSeparatingAxis = 0;
+		} else
+		{
+			bool sepB = findSeparatingAxisLocalB(	&convexShapes[shapeIndexB],&convexPolyhedronA,rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,
+																									rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,
+																									DeltaC2,
+																									vertices,uniqueEdges,faces,indices,
+																									verticesA,uniqueEdgesA,facesA,indicesA,
+																									&sepAxis,&dmin);
+
+			if (!sepB)
+			{
+				hasSeparatingAxis = 0;
+			} else
+			{
+				bool sepEE = findSeparatingAxisEdgeEdgeLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,
+																									rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,
+																									DeltaC2,
+																									verticesA,uniqueEdgesA,facesA,indicesA,
+																									vertices,uniqueEdges,faces,indices,
+																									&sepAxis,&dmin);
+	
+				if (!sepEE)
+				{
+					hasSeparatingAxis = 0;
+				} else
+				{
+					hasSeparatingAxis = 1;
+				}
+			}
+		}	
+		
+		if (hasSeparatingAxis)
+		{
+			sepAxis.w = dmin;
+			concaveSeparatingNormalsOut[pairIdx]=sepAxis;
+		} else
+		{	
+			//mark this pair as in-active
+			concavePairs[pairIdx].w = 0;
+		}
+	}
+	else
+	{	
+		//mark this pair as in-active
+		concavePairs[pairIdx].w = 0;
+	}
+}

opencl/gpu_sat/kernels/satClipHullContacts.cl

@@ -1413,6 +1413,9 @@ __kernel void   clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
 
 	if (i<numConcavePairs)
 	{
+		//magic value to detect that pair is invalid
+		if (concavePairsIn[i].w!=3)
+			return;
 
 		int bodyIndexA = concavePairsIn[i].x;
 		int bodyIndexB = concavePairsIn[i].y;

opencl/gpu_sat/kernels/satClipHullContacts.h

@@ -1415,6 +1415,9 @@ static const char* satClipKernelsCL= \
 "\n"
 "	if (i<numConcavePairs)\n"
 "	{\n"
+"		//magic value to detect that pair is invalid\n"
+"		if (concavePairsIn[i].w!=3)\n"
+"			return;\n"
 "\n"
 "		int bodyIndexA = concavePairsIn[i].x;\n"
 "		int bodyIndexB = concavePairsIn[i].y;\n"

opencl/gpu_sat/kernels/satKernels.h

@@ -941,11 +941,7 @@ static const char* satKernelsCL= \
 "																					__global btAabbCL* aabbs,\n"
 "																					__global volatile float4* separatingNormals,\n"
 "																					__global volatile int* hasSeparatingAxis,\n"
-"																					__global int4* concavePairsOut,\n"
+"																					int numPairs\n"
-"																					__global float4* concaveSeparatingNormalsOut,\n"
-"																					__global volatile int* numConcavePairsOut,\n"
-"																					int numPairs,\n"
-"																					int maxNumConcavePairsCapacity\n"
 "																					)\n"
 "{\n"
 "\n"
@@ -972,230 +968,6 @@ static const char* satKernelsCL= \
 "			return;\n"
 "		}\n"
 "		\n"
-"		if ((collidables[collidableIndexA].m_shapeType==SHAPE_CONCAVE_TRIMESH))// && (collidables[collidableIndexB].m_shapeType==SHAPE_CONVEX_HULL))\n"
-"		{\n"
-"	\n"
-"			int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
-"			int numActualConcaveConvexTests = 0;\n"
-"			\n"
-"			for (int f=0;f<numFacesA;f++)\n"
-"			{\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"
-"#if 0\n"
-"//just in case some implementation doesn't support component-wise min and max for float4\n"
-"					if (triAabb.m_min.x > 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<numVertices;i++)\n"
-"						{\n"
-"							float4 v0 = verticesA[i];\n"
-"							float4 v1 = verticesA[prevVertex];\n"
-"	                                                \n"
-"							float4 edgeNormal = normalize(cross(normal,v1-v0));\n"
-"							float c = -dot(edgeNormal,v0);\n"
-"	\n"
-"							facesA[fidx].m_numIndices = 2;\n"
-"							facesA[fidx].m_indexOffset=curUsedIndices;\n"
-"							indicesA[curUsedIndices++]=i;\n"
-"							indicesA[curUsedIndices++]=prevVertex;\n"
-"	                                                \n"
-"							facesA[fidx].m_plane.x = edgeNormal.x;\n"
-"							facesA[fidx].m_plane.y = edgeNormal.y;\n"
-"							facesA[fidx].m_plane.z = edgeNormal.z;\n"
-"							facesA[fidx].m_plane.w = c;\n"
-"							fidx++;\n"
-"							prevVertex = i;\n"
-"						}\n"
-"					}\n"
-"					convexPolyhedronA.m_numFaces = TRIANGLE_NUM_CONVEX_FACES;\n"
-"					convexPolyhedronA.m_localCenter = localCenter*(1.f/3.f);\n"
-"			\n"
-"\n"
-"					float4 posA = rigidBodies[bodyIndexA].m_pos;\n"
-"					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"
-"					bool sepA = findSeparatingAxisLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,\n"
-"																											rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,\n"
-"																											DeltaC2,\n"
-"																											verticesA,uniqueEdgesA,facesA,indicesA,\n"
-"																											vertices,uniqueEdges,faces,indices,\n"
-"																											&sepAxis,&dmin);\n"
-"					hasSeparatingAxis = 4;\n"
-"					if (!sepA)\n"
-"					{\n"
-"						hasSeparatingAxis = 0;\n"
-"					} else\n"
-"					{\n"
-"						bool sepB = findSeparatingAxisLocalB(	&convexShapes[shapeIndexB],&convexPolyhedronA,rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,\n"
-"																												rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,\n"
-"																												DeltaC2,\n"
-"																												vertices,uniqueEdges,faces,indices,\n"
-"																												verticesA,uniqueEdgesA,facesA,indicesA,\n"
-"																												&sepAxis,&dmin);\n"
-"			\n"
-"						if (!sepB)\n"
-"						{\n"
-"							hasSeparatingAxis = 0;\n"
-"						} else\n"
-"						{\n"
-"							bool sepEE = findSeparatingAxisEdgeEdgeLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,\n"
-"																												rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,\n"
-"																												DeltaC2,\n"
-"																												verticesA,uniqueEdgesA,facesA,indicesA,\n"
-"																												vertices,uniqueEdges,faces,indices,\n"
-"																												&sepAxis,&dmin);\n"
-"				\n"
-"							if (!sepEE)\n"
-"							{\n"
-"								hasSeparatingAxis = 0;\n"
-"							} else\n"
-"							{\n"
-"								hasSeparatingAxis = 1;\n"
-"							}\n"
-"						}\n"
-"					}	\n"
-"#endif\n"
-"					\n"
-"					if (hasSeparatingAxis)\n"
-"					{\n"
-"						int pairIdx = atomic_inc(numConcavePairsOut);\n"
-"						if (pairIdx<maxNumConcavePairsCapacity)\n"
-"						{\n"
-"							concavePairsOut[pairIdx].x = bodyIndexA;\n"
-"							concavePairsOut[pairIdx].y = bodyIndexB;\n"
-"							concavePairsOut[pairIdx].z = f;\n"
-"							concavePairsOut[pairIdx].w = 3;\n"
-"							sepAxis.w = dmin;\n"
-"							concaveSeparatingNormalsOut[pairIdx]=sepAxis;\n"
-"						}\n"
-"					}\n"
-"				}\n"
-"			}\n"
-"			//todo//??\n"
-"			hasSeparatingAxis[i] = 0;\n"
-"			return;\n"
-"		}		\n"
-"\n"
-"	\n"
 "\n"
 "		if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONVEX_HULL) ||(collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL))\n"
 "		{\n"
@@ -1203,10 +975,12 @@ static const char* satKernelsCL= \
 "			return;\n"
 "		}\n"
 "			\n"
+"		if ((collidables[collidableIndexA].m_shapeType==SHAPE_CONCAVE_TRIMESH))\n"
+"		{\n"
+"			hasSeparatingAxis[i] = 0;\n"
+"			return;\n"
+"		}\n"
 "\n"
-"\n"
-"\n"
-"		\n"
 "		int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
 "\n"
 "		float dmin = FLT_MAX;\n"
@@ -1265,4 +1039,238 @@ static const char* satKernelsCL= \
 "	}\n"
 "\n"
 "}\n"
+"\n"
+"\n"
+"\n"
+"\n"
+"// work-in-progress\n"
+"__kernel void   findConcaveSeparatingAxisKernel( __global int4* concavePairs,\n"
+"																					__global const BodyData* rigidBodies,\n"
+"																					__global const btCollidableGpu* collidables,\n"
+"																					__global const ConvexPolyhedronCL* convexShapes, \n"
+"																					__global const float4* vertices,\n"
+"																					__global const float4* uniqueEdges,\n"
+"																					__global const btGpuFace* faces,\n"
+"																					__global const int* indices,\n"
+"																					__global btAabbCL* aabbs,\n"
+"																					__global volatile float4* separatingNormals,\n"
+"																					__global volatile int* hasSeparatingAxis,\n"
+"																					__global float4* concaveSeparatingNormalsOut,\n"
+"																					int numConcavePairs\n"
+"																					)\n"
+"{\n"
+"\n"
+"	int i = get_global_id(0);\n"
+"	if (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<numVertices;i++)\n"
+"			{\n"
+"				float4 v0 = verticesA[i];\n"
+"				float4 v1 = verticesA[prevVertex];\n"
+"                                            \n"
+"				float4 edgeNormal = normalize(cross(normal,v1-v0));\n"
+"				float c = -dot(edgeNormal,v0);\n"
+"\n"
+"				facesA[fidx].m_numIndices = 2;\n"
+"				facesA[fidx].m_indexOffset=curUsedIndices;\n"
+"				indicesA[curUsedIndices++]=i;\n"
+"				indicesA[curUsedIndices++]=prevVertex;\n"
+"                                            \n"
+"				facesA[fidx].m_plane.x = edgeNormal.x;\n"
+"				facesA[fidx].m_plane.y = edgeNormal.y;\n"
+"				facesA[fidx].m_plane.z = edgeNormal.z;\n"
+"				facesA[fidx].m_plane.w = c;\n"
+"				fidx++;\n"
+"				prevVertex = i;\n"
+"			}\n"
+"		}\n"
+"		convexPolyhedronA.m_numFaces = TRIANGLE_NUM_CONVEX_FACES;\n"
+"		convexPolyhedronA.m_localCenter = localCenter*(1.f/3.f);\n"
+"\n"
+"\n"
+"		float4 posA = rigidBodies[bodyIndexA].m_pos;\n"
+"		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"
+"		bool sepA = findSeparatingAxisLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,\n"
+"																								rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,\n"
+"																								DeltaC2,\n"
+"																								verticesA,uniqueEdgesA,facesA,indicesA,\n"
+"																								vertices,uniqueEdges,faces,indices,\n"
+"																								&sepAxis,&dmin);\n"
+"		hasSeparatingAxis = 4;\n"
+"		if (!sepA)\n"
+"		{\n"
+"			hasSeparatingAxis = 0;\n"
+"		} else\n"
+"		{\n"
+"			bool sepB = findSeparatingAxisLocalB(	&convexShapes[shapeIndexB],&convexPolyhedronA,rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,\n"
+"																									rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,\n"
+"																									DeltaC2,\n"
+"																									vertices,uniqueEdges,faces,indices,\n"
+"																									verticesA,uniqueEdgesA,facesA,indicesA,\n"
+"																									&sepAxis,&dmin);\n"
+"\n"
+"			if (!sepB)\n"
+"			{\n"
+"				hasSeparatingAxis = 0;\n"
+"			} else\n"
+"			{\n"
+"				bool sepEE = findSeparatingAxisEdgeEdgeLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,\n"
+"																									rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,\n"
+"																									DeltaC2,\n"
+"																									verticesA,uniqueEdgesA,facesA,indicesA,\n"
+"																									vertices,uniqueEdges,faces,indices,\n"
+"																									&sepAxis,&dmin);\n"
+"	\n"
+"				if (!sepEE)\n"
+"				{\n"
+"					hasSeparatingAxis = 0;\n"
+"				} else\n"
+"				{\n"
+"					hasSeparatingAxis = 1;\n"
+"				}\n"
+"			}\n"
+"		}	\n"
+"		\n"
+"		if (hasSeparatingAxis)\n"
+"		{\n"
+"			sepAxis.w = dmin;\n"
+"			concaveSeparatingNormalsOut[pairIdx]=sepAxis;\n"
+"		} else\n"
+"		{	\n"
+"			//mark this pair as in-active\n"
+"			concavePairs[pairIdx].w = 0;\n"
+"		}\n"
+"	}\n"
+"	else\n"
+"	{	\n"
+"		//mark this pair as in-active\n"
+"		concavePairs[pairIdx].w = 0;\n"
+"	}\n"
+"}\n"
+"\n"
 ;