From c8fcf779bbbe7aeb75696dc58728e9b453e2642b Mon Sep 17 00:00:00 2001
From: erwin coumans <erwin.coumans@gmail.com>
Date: Fri, 5 Apr 2013 18:29:58 -0700
Subject: [PATCH] add support for compound hulls against concave triangle mesh

---
 demo/gpudemo/GpuDemo.h                        |   6 +-
 demo/gpudemo/main_opengl3core.cpp             |   6 +-
 demo/gpudemo/rigidbody/ConcaveScene.cpp       | 129 +++++++++++++++++-
 demo/gpudemo/rigidbody/ConcaveScene.h         |  25 ++++
 opencl/gpu_rigidbody/host/btConfig.h          |   2 +-
 opencl/gpu_sat/host/ConvexHullContact.cpp     |  17 ++-
 opencl/gpu_sat/kernels/bvhTraversal.cl        |  29 +++-
 opencl/gpu_sat/kernels/bvhTraversal.h         |  29 +++-
 opencl/gpu_sat/kernels/sat.cl                 |  75 ++++++----
 opencl/gpu_sat/kernels/satClipHullContacts.cl |  25 +++-
 opencl/gpu_sat/kernels/satClipHullContacts.h  |  25 +++-
 opencl/gpu_sat/kernels/satKernels.h           |  75 ++++++----
 12 files changed, 361 insertions(+), 82 deletions(-)

diff --git a/demo/gpudemo/GpuDemo.h b/demo/gpudemo/GpuDemo.h
index 77af3c89d..c62028709 100644
--- a/demo/gpudemo/GpuDemo.h
+++ b/demo/gpudemo/GpuDemo.h
@@ -38,9 +38,9 @@ public:
                     :useOpenCL(true),
                     preferredOpenCLPlatformIndex(-1),
                     preferredOpenCLDeviceIndex(-1),
-					arraySizeX(41),
-		arraySizeY(41),
-		arraySizeZ(41),
+					arraySizeX(11),
+		arraySizeY(11),
+		arraySizeZ(11),
 		m_useConcaveMesh(false),
 		gapX(14.3),
 		gapY(14.0),
diff --git a/demo/gpudemo/main_opengl3core.cpp b/demo/gpudemo/main_opengl3core.cpp
index 08c5fadd3..6f6e70834 100644
--- a/demo/gpudemo/main_opengl3core.cpp
+++ b/demo/gpudemo/main_opengl3core.cpp
@@ -65,14 +65,16 @@ btAlignedObjectArray<const char*> demoNames;
 int selectedDemo = 0;
 GpuDemo::CreateFunc* allDemos[]=
 {
+	ConcaveCompoundScene::MyCreateFunc,
 
+	ConcaveSphereScene::MyCreateFunc,
+
+	ConcaveScene::MyCreateFunc,
 	
 	GpuSphereScene::MyCreateFunc,
 	
 	
 	GpuConvexScene::MyCreateFunc,
-	ConcaveScene::MyCreateFunc,
-	ConcaveCompoundScene::MyCreateFunc,
 	
 	GpuCompoundScene::MyCreateFunc,
 	PairBench::MyCreateFunc,	
diff --git a/demo/gpudemo/rigidbody/ConcaveScene.cpp b/demo/gpudemo/rigidbody/ConcaveScene.cpp
index 17962a411..7e1f03bac 100644
--- a/demo/gpudemo/rigidbody/ConcaveScene.cpp
+++ b/demo/gpudemo/rigidbody/ConcaveScene.cpp
@@ -17,7 +17,7 @@
 #include "gpu_rigidbody/host/btConfig.h"
 #include "GpuRigidBodyDemoInternalData.h"
 #include"../../ObjLoader/objLoader.h"
-
+#include "BulletCommon/btTransform.h"
 
 #include "OpenGLWindow/GLInstanceGraphicsShape.h"
 #define CONCAVE_GAPX 16
@@ -328,6 +328,133 @@ void ConcaveCompoundScene::setupScene(const ConstructionInfo& ci)
 }
 
 void ConcaveCompoundScene::createDynamicObjects(const ConstructionInfo& ci)
+{
+	
+		int strideInBytes = 9*sizeof(float);
+	int numVertices = sizeof(cube_vertices)/strideInBytes;
+	int numIndices = sizeof(cube_indices)/sizeof(int);
+
+	btAlignedObjectArray<GLInstanceVertex> vertexArray;
+	btAlignedObjectArray<int> indexArray;
+	
+
+	//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
+	int group=1;
+	int mask=1;
+	int index=0;
+	float scaling[4] = {1,1,1,1};
+	int colIndex = 0;
+
+	GLInstanceVertex* cubeVerts = (GLInstanceVertex*)&cube_vertices[0];
+	int stride2 = sizeof(GLInstanceVertex);
+	btAssert(stride2 == strideInBytes);
+
+	{
+		int childColIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
+		
+
+		btVector3 childPositions[3] = {
+			btVector3(0,-2,0),
+			btVector3(0,0,0),
+			btVector3(0,2,0)
+		};
+
+		
+		btAlignedObjectArray<btGpuChildShape> childShapes;
+		int numChildShapes = 3;
+		for (int i=0;i<numChildShapes;i++)
+		{
+			//for now, only support polyhedral child shapes
+			btGpuChildShape child;
+			child.m_shapeIndex = childColIndex;
+			btVector3 pos = childPositions[i];
+			btQuaternion orn(0,0,0,1);
+			for (int v=0;v<4;v++)
+			{
+				child.m_childPosition[v] = pos[v];
+				child.m_childOrientation[v] = orn[v];
+			}
+			childShapes.push_back(child);
+			btTransform tr;
+			tr.setIdentity();
+			tr.setOrigin(pos);
+			tr.setRotation(orn);
+
+			int baseIndex = vertexArray.size();
+			for (int j=0;j<numIndices;j++)
+				indexArray.push_back(cube_indices[j]+baseIndex);
+
+			//add transformed graphics vertices and indices
+			for (int v=0;v<numVertices;v++)
+			{
+				GLInstanceVertex vert = cubeVerts[v];
+				btVector3 vertPos(vert.xyzw[0],vert.xyzw[1],vert.xyzw[2]);
+				btVector3 newPos = tr*vertPos;
+				vert.xyzw[0] = newPos[0];
+				vert.xyzw[1] = newPos[1];
+				vert.xyzw[2] = newPos[2];
+				vert.xyzw[3] = 0.f;
+				vertexArray.push_back(vert);
+			}
+
+		}
+		colIndex= m_data->m_np->registerCompoundShape(&childShapes);
+		
+	}
+
+	//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
+	int shapeId = ci.m_instancingRenderer->registerShape(&vertexArray[0].xyzw[0],vertexArray.size(),&indexArray[0],indexArray.size());
+
+	btVector4 colors[4] = 
+	{
+		btVector4(1,0,0,1),
+		btVector4(0,1,0,1),
+		btVector4(0,0,1,1),
+		btVector4(0,1,1,1),
+	};
+		
+	int curColor = 0;
+	for (int i=0;i<ci.arraySizeX;i++)
+	{
+		for (int j=0;j<ci.arraySizeY;j++)
+		{
+			for (int k=0;k<ci.arraySizeZ;k++)
+
+			{
+				float mass = 1;//j==0? 0.f : 1.f;
+
+				btVector3 position(i*ci.gapX,10+j*ci.gapY,k*ci.gapZ);
+				//btQuaternion orn(0,0,0,1);
+				btQuaternion orn(btVector3(1,0,0),0.7);
+				
+				btVector4 color = colors[curColor];
+				curColor++;
+				curColor&=3;
+				btVector4 scaling(1,1,1,1);
+				int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
+				int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(mass,position,orn,colIndex,index);
+				
+				index++;
+			}
+		}
+	}
+}
+
+
+void ConcaveSphereScene::setupScene(const ConstructionInfo& ci)
+{
+	ConcaveScene::setupScene(ci);
+
+	float camPos[4]={0,50,0,0};//65.5,4.5,65.5,0};
+	//float camPos[4]={1,12.5,1.5,0};
+	m_instancingRenderer->setCameraPitch(45);
+	m_instancingRenderer->setCameraTargetPosition(camPos);
+	m_instancingRenderer->setCameraDistance(40);
+
+
+}
+
+void ConcaveSphereScene::createDynamicObjects(const ConstructionInfo& ci)
 {
 		btVector4 colors[4] = 
 	{
diff --git a/demo/gpudemo/rigidbody/ConcaveScene.h b/demo/gpudemo/rigidbody/ConcaveScene.h
index 1fab1f74f..617cb5538 100644
--- a/demo/gpudemo/rigidbody/ConcaveScene.h
+++ b/demo/gpudemo/rigidbody/ConcaveScene.h
@@ -29,6 +29,31 @@ public:
 	
 };
 
+
+class ConcaveSphereScene : public ConcaveScene
+{
+public:
+
+	ConcaveSphereScene(){}
+	virtual ~ConcaveSphereScene(){}
+	virtual const char* getName()
+	{
+		return "ConcaveSphere";
+	}
+
+	static GpuDemo* MyCreateFunc()
+	{
+		GpuDemo* demo = new ConcaveSphereScene;
+		return demo;
+	}
+
+	virtual void setupScene(const ConstructionInfo& ci);
+
+	virtual void createDynamicObjects(const ConstructionInfo& ci);
+
+};
+
+
 class ConcaveCompoundScene : public ConcaveScene
 {
 public:
diff --git a/opencl/gpu_rigidbody/host/btConfig.h b/opencl/gpu_rigidbody/host/btConfig.h
index 25ca1ad89..77ddec96b 100644
--- a/opencl/gpu_rigidbody/host/btConfig.h
+++ b/opencl/gpu_rigidbody/host/btConfig.h
@@ -19,7 +19,7 @@ struct	btConfig
 	int m_maxTriConvexPairCapacity;
 
 	btConfig()
-		:m_maxConvexBodies(128*1024),
+		:m_maxConvexBodies(32*1024),
 		m_maxConvexShapes(8192),
 		m_maxVerticesPerFace(64),
 		m_maxFacesPerShape(64),
diff --git a/opencl/gpu_sat/host/ConvexHullContact.cpp b/opencl/gpu_sat/host/ConvexHullContact.cpp
index 679b118f8..ca5d55610 100644
--- a/opencl/gpu_sat/host/ConvexHullContact.cpp
+++ b/opencl/gpu_sat/host/ConvexHullContact.cpp
@@ -58,7 +58,7 @@ m_totalContactsOut(m_context, m_queue)
 	if (1)
 	{
 		const char* src = satKernelsCL;
-		cl_program satProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,src,&errNum,"","opencl/gpu_sat/kernels/sat.cl");
+		cl_program satProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,src,&errNum,"","opencl/gpu_sat/kernels/sat.cl",true);
 		btAssert(errNum==CL_SUCCESS);
 
 		m_findSeparatingAxisKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,src, "findSeparatingAxisKernel",&errNum,satProg );
@@ -80,7 +80,7 @@ m_totalContactsOut(m_context, m_queue)
 	if (1)
 	{
 		const char* srcClip = satClipKernelsCL;
-		cl_program satClipContactsProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,srcClip,&errNum,"","opencl/gpu_sat/kernels/satClipHullContacts.cl");
+		cl_program satClipContactsProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,srcClip,&errNum,"","opencl/gpu_sat/kernels/satClipHullContacts.cl",true);
 		btAssert(errNum==CL_SUCCESS);
 
 		m_clipHullHullKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,srcClip, "clipHullHullKernel",&errNum,satClipContactsProg);
@@ -120,7 +120,7 @@ m_totalContactsOut(m_context, m_queue)
 	 if (1)
 	{
 		const char* srcBvh = bvhTraversalKernelCL;
-		cl_program bvhTraversalProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,srcBvh,&errNum,"","opencl/gpu_sat/kernels/bvhTraversal.cl");
+		cl_program bvhTraversalProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,srcBvh,&errNum,"","opencl/gpu_sat/kernels/bvhTraversal.cl",true);
 		//cl_program bvhTraversalProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,0,&errNum,"","opencl/gpu_sat/kernels/bvhTraversal.cl", true);
 		btAssert(errNum==CL_SUCCESS);
 
@@ -706,9 +706,8 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 								btBufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
 								btBufferInfoCL( gpuFaces.getBufferCL(),true),
 								btBufferInfoCL( gpuIndices.getBufferCL(),true),
+								btBufferInfoCL( gpuChildShapes.getBufferCL(),true),
 								btBufferInfoCL( clAabbsWS.getBufferCL(),true),
-								btBufferInfoCL( sepNormals.getBufferCL()),
-								btBufferInfoCL( hasSeparatingNormals.getBufferCL()),
 								btBufferInfoCL( concaveSepNormals.getBufferCL())
 							};
 
@@ -725,7 +724,9 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 				}
 			}
 			
+			numCompoundPairs = numCompoundPairsOut.at(0);
 
+			if (1)
 			{
 				BT_PROFILE("findCompoundPairsKernel");
 				btBufferInfoCL bInfo[] = 
@@ -896,6 +897,10 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 
 		if (numConcavePairs)
 		{
+//			printf("numConcavePairs = %d\n", numConcavePairs);
+	//		nContacts = m_totalContactsOut.at(0);
+		//	printf("nContacts before = %d\n", nContacts);
+
 			BT_PROFILE("clipHullHullConcaveConvexKernel");
 			nContacts = m_totalContactsOut.at(0);
 			btBufferInfoCL bInfo[] = { 
@@ -907,6 +912,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 				btBufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
 				btBufferInfoCL( gpuFaces.getBufferCL(),true),
 				btBufferInfoCL( gpuIndices.getBufferCL(),true),
+				btBufferInfoCL( gpuChildShapes.getBufferCL(),true),
 				btBufferInfoCL( concaveSepNormals.getBufferCL()),
 				btBufferInfoCL( contactOut->getBufferCL()),
 				btBufferInfoCL( m_totalContactsOut.getBufferCL())	
@@ -918,6 +924,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 			launcher.launch1D( num);
 			clFinish(m_queue);
 			nContacts = m_totalContactsOut.at(0);
+			//printf("nContacts after = %d\n", nContacts);
 		}
 
 		
diff --git a/opencl/gpu_sat/kernels/bvhTraversal.cl b/opencl/gpu_sat/kernels/bvhTraversal.cl
index a0e3e07c3..33795628a 100644
--- a/opencl/gpu_sat/kernels/bvhTraversal.cl
+++ b/opencl/gpu_sat/kernels/bvhTraversal.cl
@@ -215,7 +215,9 @@ __kernel void   bvhTraversalKernel( __global const int2* pairs,
 	int shapeTypeB = collidables[collidableIndexB].m_shapeType;
 		
 	if (shapeTypeB!=SHAPE_CONVEX_HULL &&
-		shapeTypeB!=SHAPE_SPHERE	)
+		shapeTypeB!=SHAPE_SPHERE	&&
+		shapeTypeB!=SHAPE_COMPOUND_OF_CONVEX_HULLS
+		)
 		return;
 
 	
@@ -247,12 +249,27 @@ __kernel void   bvhTraversalKernel( __global const int2* pairs,
 					if (isLeafNode)
 					{
 						int triangleIndex = getTriangleIndex(&rootNode);
-							
-						int pairIdx = atomic_inc(numConcavePairsOut);
-						if (pairIdx<maxNumConcavePairsCapacity)
+						if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS)
 						{
-							int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,3);
-							concavePairsOut[pairIdx] = newPair;
+								int numChildrenB = collidables[collidableIndexB].m_numChildShapes;
+								int pairIdx = atomic_add(numConcavePairsOut,numChildrenB);
+								for (int b=0;b<numChildrenB;b++)
+								{
+									if ((pairIdx+b)<maxNumConcavePairsCapacity)
+									{
+										int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+b;
+										int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,childShapeIndexB);
+										concavePairsOut[pairIdx+b] = newPair;
+									}
+								}
+						} else
+						{
+							int pairIdx = atomic_inc(numConcavePairsOut);
+							if (pairIdx<maxNumConcavePairsCapacity)
+							{
+								int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,0);
+								concavePairsOut[pairIdx] = newPair;
+							}
 						}
 					} 
 					curIndex++;
diff --git a/opencl/gpu_sat/kernels/bvhTraversal.h b/opencl/gpu_sat/kernels/bvhTraversal.h
index 2d0825158..93fb018e4 100644
--- a/opencl/gpu_sat/kernels/bvhTraversal.h
+++ b/opencl/gpu_sat/kernels/bvhTraversal.h
@@ -217,7 +217,9 @@ static const char* bvhTraversalKernelCL= \
 "	int shapeTypeB = collidables[collidableIndexB].m_shapeType;\n"
 "		\n"
 "	if (shapeTypeB!=SHAPE_CONVEX_HULL &&\n"
-"		shapeTypeB!=SHAPE_SPHERE	)\n"
+"		shapeTypeB!=SHAPE_SPHERE	&&\n"
+"		shapeTypeB!=SHAPE_COMPOUND_OF_CONVEX_HULLS\n"
+"		)\n"
 "		return;\n"
 "\n"
 "	\n"
@@ -249,12 +251,27 @@ static const char* bvhTraversalKernelCL= \
 "					if (isLeafNode)\n"
 "					{\n"
 "						int triangleIndex = getTriangleIndex(&rootNode);\n"
-"							\n"
-"						int pairIdx = atomic_inc(numConcavePairsOut);\n"
-"						if (pairIdx<maxNumConcavePairsCapacity)\n"
+"						if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
 "						{\n"
-"							int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,3);\n"
-"							concavePairsOut[pairIdx] = newPair;\n"
+"								int numChildrenB = collidables[collidableIndexB].m_numChildShapes;\n"
+"								int pairIdx = atomic_add(numConcavePairsOut,numChildrenB);\n"
+"								for (int b=0;b<numChildrenB;b++)\n"
+"								{\n"
+"									if ((pairIdx+b)<maxNumConcavePairsCapacity)\n"
+"									{\n"
+"										int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+b;\n"
+"										int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,childShapeIndexB);\n"
+"										concavePairsOut[pairIdx+b] = newPair;\n"
+"									}\n"
+"								}\n"
+"						} else\n"
+"						{\n"
+"							int pairIdx = atomic_inc(numConcavePairsOut);\n"
+"							if (pairIdx<maxNumConcavePairsCapacity)\n"
+"							{\n"
+"								int4 newPair = (int4)(bodyIndexA,bodyIndexB,triangleIndex,0);\n"
+"								concavePairsOut[pairIdx] = newPair;\n"
+"							}\n"
 "						}\n"
 "					} \n"
 "					curIndex++;\n"
diff --git a/opencl/gpu_sat/kernels/sat.cl b/opencl/gpu_sat/kernels/sat.cl
index 465cdd86e..e06fc85d6 100644
--- a/opencl/gpu_sat/kernels/sat.cl
+++ b/opencl/gpu_sat/kernels/sat.cl
@@ -886,7 +886,8 @@ __kernel void   findCompoundPairsKernel( __global const int2* pairs,
 				}//for (int b=0;b<numChildrenB;b++)	
 				return;
 			}//if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)
-			if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)
+			if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONCAVE_TRIMESH) 
+				&& (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS))
 			{
 				int numChildrenB = collidables[collidableIndexB].m_numChildShapes;
 				for (int b=0;b<numChildrenB;b++)
@@ -1059,9 +1060,8 @@ __kernel void   findConcaveSeparatingAxisKernel( __global int4* concavePairs,
 																					__global const float4* uniqueEdges,
 																					__global const btGpuFace* faces,
 																					__global const int* indices,
+																					__global const btGpuChildShape* gpuChildShapes,
 																					__global btAabbCL* aabbs,
-																					__global volatile float4* separatingNormals,
-																					__global volatile int* hasSeparatingAxis,
 																					__global float4* concaveSeparatingNormalsOut,
 																					int numConcavePairs
 																					)
@@ -1081,9 +1081,10 @@ __kernel void   findConcaveSeparatingAxisKernel( __global int4* concavePairs,
 	int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;
 	int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;
 
-	if (collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL)
+	if (collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL&&
+		collidables[collidableIndexB].m_shapeType!=SHAPE_COMPOUND_OF_CONVEX_HULLS)
 	{
-		concavePairs[pairIdx].w = 0;
+		concavePairs[pairIdx].w = -1;
 		return;
 	}
 
@@ -1229,36 +1230,60 @@ __kernel void   findConcaveSeparatingAxisKernel( __global int4* concavePairs,
 		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);
+
+
+		///////////////////
+		///compound shape support
+
+		if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)
+		{
+			int compoundChild = concavePairs[pairIdx].w;
+			int childShapeIndexB = compoundChild;//collidables[collidableIndexB].m_shapeIndex+compoundChild;
+			int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;
+			float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;
+			float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;
+			float4 newPosB = transform(&childPosB,&posB,&ornB);
+			float4 newOrnB = qtMul(ornB,childOrnB);
+			posB = newPosB;
+			ornB = newOrnB;
+			shapeIndexB = collidables[childColIndexB].m_shapeIndex;
+		}
+		//////////////////
+
+
+		bool sepA = findSeparatingAxisLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],
+												posA,ornA,
+												posB,ornB,
+												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);
+			bool sepB = findSeparatingAxisLocalB(	&convexShapes[shapeIndexB],&convexPolyhedronA,
+												posB,ornB,
+												posA,ornA,
+												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);
+				bool sepEE = findSeparatingAxisEdgeEdgeLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],
+															posA,ornA,
+															posB,ornB,
+															DeltaC2,
+															verticesA,uniqueEdgesA,facesA,indicesA,
+															vertices,uniqueEdges,faces,indices,
+															&sepAxis,&dmin);
 	
 				if (!sepEE)
 				{
@@ -1277,12 +1302,12 @@ __kernel void   findConcaveSeparatingAxisKernel( __global int4* concavePairs,
 		} else
 		{	
 			//mark this pair as in-active
-			concavePairs[pairIdx].w = 0;
+			concavePairs[pairIdx].w = -1;
 		}
 	}
 	else
 	{	
 		//mark this pair as in-active
-		concavePairs[pairIdx].w = 0;
+		concavePairs[pairIdx].w = -1;
 	}
 }
diff --git a/opencl/gpu_sat/kernels/satClipHullContacts.cl b/opencl/gpu_sat/kernels/satClipHullContacts.cl
index d92898a53..4d3f5140e 100644
--- a/opencl/gpu_sat/kernels/satClipHullContacts.cl
+++ b/opencl/gpu_sat/kernels/satClipHullContacts.cl
@@ -1256,6 +1256,7 @@ __kernel void   clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
 																					__global const float4* uniqueEdges,
 																					__global const btGpuFace* faces,
 																					__global const int* indices,
+																					__global const btGpuChildShape* gpuChildShapes,
 																					__global const float4* separatingNormals,
 																					__global Contact4* restrict globalContactsOut,
 																					counter32_t nGlobalContactsOut,
@@ -1277,8 +1278,8 @@ __kernel void   clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
 
 	if (i<numConcavePairs)
 	{
-		//magic value to detect that pair is invalid
-		if (concavePairsIn[i].w!=3)
+		//negative value means that the pair is invalid
+		if (concavePairsIn[i].w<0)
 			return;
 
 		int bodyIndexA = concavePairsIn[i].x;
@@ -1413,6 +1414,22 @@ __kernel void   clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
 
 		float4 sepAxis = separatingNormals[i];
 		
+		int shapeTypeB = collidables[collidableIndexB].m_shapeType;
+		if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS)
+		{
+			///////////////////
+			///compound shape support
+			int compoundChild = concavePairsIn[pairIndex].w;
+			int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+compoundChild;
+			int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;
+			float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;
+			float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;
+			float4 newPosB = transform(&childPosB,&posB,&ornB);
+			float4 newOrnB = qtMul(ornB,childOrnB);
+			posB = newPosB;
+			ornB = newOrnB;
+			shapeIndexB = collidables[childColIndexB].m_shapeIndex;
+		}
 		
 		////////////////////////////////////////
 		
@@ -1420,8 +1437,8 @@ __kernel void   clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
 		
 		int numLocalContactsOut = clipHullAgainstHullLocalA(sepAxis,
 														&convexPolyhedronA, &convexShapes[shapeIndexB],
-														rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,
-													  rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,
+														posA,ornA,
+													  posB,ornB,
 													  worldVertsB1,worldVertsB2,capacityWorldVerts,
 														minDist, maxDist,
 														&verticesA,&facesA,&indicesA,
diff --git a/opencl/gpu_sat/kernels/satClipHullContacts.h b/opencl/gpu_sat/kernels/satClipHullContacts.h
index fc9e50d1b..c8be64714 100644
--- a/opencl/gpu_sat/kernels/satClipHullContacts.h
+++ b/opencl/gpu_sat/kernels/satClipHullContacts.h
@@ -1258,6 +1258,7 @@ static const char* satClipKernelsCL= \
 "																					__global const float4* uniqueEdges,\n"
 "																					__global const btGpuFace* faces,\n"
 "																					__global const int* indices,\n"
+"																					__global const btGpuChildShape* gpuChildShapes,\n"
 "																					__global const float4* separatingNormals,\n"
 "																					__global Contact4* restrict globalContactsOut,\n"
 "																					counter32_t nGlobalContactsOut,\n"
@@ -1279,8 +1280,8 @@ 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"
+"		//negative value means that the pair is invalid\n"
+"		if (concavePairsIn[i].w<0)\n"
 "			return;\n"
 "\n"
 "		int bodyIndexA = concavePairsIn[i].x;\n"
@@ -1415,6 +1416,22 @@ static const char* satClipKernelsCL= \
 "\n"
 "		float4 sepAxis = separatingNormals[i];\n"
 "		\n"
+"		int shapeTypeB = collidables[collidableIndexB].m_shapeType;\n"
+"		if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
+"		{\n"
+"			///////////////////\n"
+"			///compound shape support\n"
+"			int compoundChild = concavePairsIn[pairIndex].w;\n"
+"			int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+compoundChild;\n"
+"			int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;\n"
+"			float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;\n"
+"			float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;\n"
+"			float4 newPosB = transform(&childPosB,&posB,&ornB);\n"
+"			float4 newOrnB = qtMul(ornB,childOrnB);\n"
+"			posB = newPosB;\n"
+"			ornB = newOrnB;\n"
+"			shapeIndexB = collidables[childColIndexB].m_shapeIndex;\n"
+"		}\n"
 "		\n"
 "		////////////////////////////////////////\n"
 "		\n"
@@ -1422,8 +1439,8 @@ static const char* satClipKernelsCL= \
 "		\n"
 "		int numLocalContactsOut = clipHullAgainstHullLocalA(sepAxis,\n"
 "														&convexPolyhedronA, &convexShapes[shapeIndexB],\n"
-"														rigidBodies[bodyIndexA].m_pos,rigidBodies[bodyIndexA].m_quat,\n"
-"													  rigidBodies[bodyIndexB].m_pos,rigidBodies[bodyIndexB].m_quat,\n"
+"														posA,ornA,\n"
+"													  posB,ornB,\n"
 "													  worldVertsB1,worldVertsB2,capacityWorldVerts,\n"
 "														minDist, maxDist,\n"
 "														&verticesA,&facesA,&indicesA,\n"
diff --git a/opencl/gpu_sat/kernels/satKernels.h b/opencl/gpu_sat/kernels/satKernels.h
index 6c1cd700f..79bf3c5cd 100644
--- a/opencl/gpu_sat/kernels/satKernels.h
+++ b/opencl/gpu_sat/kernels/satKernels.h
@@ -888,7 +888,8 @@ static const char* satKernelsCL= \
 "				}//for (int b=0;b<numChildrenB;b++)	\n"
 "				return;\n"
 "			}//if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
-"			if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
+"			if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONCAVE_TRIMESH) \n"
+"				&& (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS))\n"
 "			{\n"
 "				int numChildrenB = collidables[collidableIndexB].m_numChildShapes;\n"
 "				for (int b=0;b<numChildrenB;b++)\n"
@@ -1061,9 +1062,8 @@ static const char* satKernelsCL= \
 "																					__global const float4* uniqueEdges,\n"
 "																					__global const btGpuFace* faces,\n"
 "																					__global const int* indices,\n"
+"																					__global const btGpuChildShape* gpuChildShapes,\n"
 "																					__global btAabbCL* aabbs,\n"
-"																					__global volatile float4* separatingNormals,\n"
-"																					__global volatile int* hasSeparatingAxis,\n"
 "																					__global float4* concaveSeparatingNormalsOut,\n"
 "																					int numConcavePairs\n"
 "																					)\n"
@@ -1083,9 +1083,10 @@ static const char* satKernelsCL= \
 "	int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;\n"
 "	int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;\n"
 "\n"
-"	if (collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL)\n"
+"	if (collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL&&\n"
+"		collidables[collidableIndexB].m_shapeType!=SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
 "	{\n"
-"		concavePairs[pairIdx].w = 0;\n"
+"		concavePairs[pairIdx].w = -1;\n"
 "		return;\n"
 "	}\n"
 "\n"
@@ -1231,36 +1232,60 @@ static const char* satKernelsCL= \
 "		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"
+"\n"
+"\n"
+"		///////////////////\n"
+"		///compound shape support\n"
+"\n"
+"		if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
+"		{\n"
+"			int compoundChild = concavePairs[pairIdx].w;\n"
+"			int childShapeIndexB = compoundChild;//collidables[collidableIndexB].m_shapeIndex+compoundChild;\n"
+"			int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;\n"
+"			float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;\n"
+"			float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;\n"
+"			float4 newPosB = transform(&childPosB,&posB,&ornB);\n"
+"			float4 newOrnB = qtMul(ornB,childOrnB);\n"
+"			posB = newPosB;\n"
+"			ornB = newOrnB;\n"
+"			shapeIndexB = collidables[childColIndexB].m_shapeIndex;\n"
+"		}\n"
+"		//////////////////\n"
+"\n"
+"\n"
+"		bool sepA = findSeparatingAxisLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],\n"
+"												posA,ornA,\n"
+"												posB,ornB,\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"
+"			bool sepB = findSeparatingAxisLocalB(	&convexShapes[shapeIndexB],&convexPolyhedronA,\n"
+"												posB,ornB,\n"
+"												posA,ornA,\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"
+"				bool sepEE = findSeparatingAxisEdgeEdgeLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],\n"
+"															posA,ornA,\n"
+"															posB,ornB,\n"
+"															DeltaC2,\n"
+"															verticesA,uniqueEdgesA,facesA,indicesA,\n"
+"															vertices,uniqueEdges,faces,indices,\n"
+"															&sepAxis,&dmin);\n"
 "	\n"
 "				if (!sepEE)\n"
 "				{\n"
@@ -1279,13 +1304,13 @@ static const char* satKernelsCL= \
 "		} else\n"
 "		{	\n"
 "			//mark this pair as in-active\n"
-"			concavePairs[pairIdx].w = 0;\n"
+"			concavePairs[pairIdx].w = -1;\n"
 "		}\n"
 "	}\n"
 "	else\n"
 "	{	\n"
 "		//mark this pair as in-active\n"
-"		concavePairs[pairIdx].w = 0;\n"
+"		concavePairs[pairIdx].w = -1;\n"
 "	}\n"
 "}\n"
 "\n"