diff --git a/src/Bullet3Collision/NarrowPhaseCollision/shared/b3FindConcaveSatAxis.h b/src/Bullet3Collision/NarrowPhaseCollision/shared/b3FindConcaveSatAxis.h
index 6a080ea98..0026df238 100644
--- a/src/Bullet3Collision/NarrowPhaseCollision/shared/b3FindConcaveSatAxis.h
+++ b/src/Bullet3Collision/NarrowPhaseCollision/shared/b3FindConcaveSatAxis.h
@@ -175,8 +175,8 @@ bool b3FindSeparatingAxisEdgeEdge(	const b3ConvexPolyhedronData* hullA, __global
 					b3Project(hullB,posB,ornB,&crossje,verticesB, &Min1, &Max1);
 				
 					if(Max0<Min1 || Max1<Min0)
-						result = false;
-				
+						return false;
+                    
 					float d0 = Max0 - Min1;
 					float d1 = Max1 - Min0;
 					dist = d0<d1 ? d0:d1;
diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.cpp b/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.cpp
index e8ea5cb23..e89b52853 100644
--- a/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.cpp
+++ b/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.cpp
@@ -14,7 +14,8 @@ subject to the following restrictions:
 */
 
 bool findSeparatingAxisOnGpu = true;
-bool splitSearchSepAxis = false;//true;
+bool splitSearchSepAxisConcave = false;
+bool splitSearchSepAxisConvex = true;
 
 bool bvhTraversalKernelGPU = true;
 bool findConcaveSeparatingAxisKernelGPU = true;
@@ -134,6 +135,17 @@ m_dmins(m_context,m_queue)
 		m_findConcaveSeparatingAxisKernel = b3OpenCLUtils::compileCLKernelFromString(m_context, m_device,src, "findConcaveSeparatingAxisKernel",&errNum,satProg );
 		b3Assert(m_findConcaveSeparatingAxisKernel);
 		b3Assert(errNum==CL_SUCCESS);
+        
+        m_findConcaveSeparatingAxisVertexFaceKernel = b3OpenCLUtils::compileCLKernelFromString(m_context, m_device,src, "findConcaveSeparatingAxisVertexFaceKernel",&errNum,satProg );
+		b3Assert(m_findConcaveSeparatingAxisVertexFaceKernel);
+		b3Assert(errNum==CL_SUCCESS);
+        
+        m_findConcaveSeparatingAxisEdgeEdgeKernel = b3OpenCLUtils::compileCLKernelFromString(m_context, m_device,src, "findConcaveSeparatingAxisEdgeEdgeKernel",&errNum,satProg );
+		b3Assert(m_findConcaveSeparatingAxisEdgeEdgeKernel);
+		b3Assert(errNum==CL_SUCCESS);
+        
+     
+        
 		
 		m_findCompoundPairsKernel = b3OpenCLUtils::compileCLKernelFromString(m_context, m_device,src, "findCompoundPairsKernel",&errNum,satProg );
 		b3Assert(m_findCompoundPairsKernel);
@@ -234,6 +246,13 @@ GpuSatCollision::~GpuSatCollision()
 	if (m_findSeparatingAxisKernel)
 		clReleaseKernel(m_findSeparatingAxisKernel);
 
+    if (m_findConcaveSeparatingAxisVertexFaceKernel)
+        clReleaseKernel(m_findConcaveSeparatingAxisVertexFaceKernel);
+
+    
+    if (m_findConcaveSeparatingAxisEdgeEdgeKernel)
+        clReleaseKernel(m_findConcaveSeparatingAxisEdgeEdgeKernel);
+    
 	if (m_findConcaveSeparatingAxisKernel)
 		clReleaseKernel(m_findConcaveSeparatingAxisKernel);
 
@@ -3039,7 +3058,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
 		if (findSeparatingAxisOnGpu)
 		{
 			m_dmins.resize(nPairs);
-			if (splitSearchSepAxis)
+			if (splitSearchSepAxisConvex)
 			{
 				{
 				B3_PROFILE("findSeparatingAxisVertexFaceKernel");
@@ -3119,211 +3138,338 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
 			}
 			
 			
-			numCompoundPairs = m_numCompoundPairsOut.at(0);
-			bool useGpuFindCompoundPairs=true;
-			if (useGpuFindCompoundPairs)
-			{
-				B3_PROFILE("findCompoundPairsKernel");
-				b3BufferInfoCL bInfo[] = 
-				{ 
-					b3BufferInfoCL( pairs->getBufferCL(), true ), 
-					b3BufferInfoCL( bodyBuf->getBufferCL(),true), 
-					b3BufferInfoCL( gpuCollidables.getBufferCL(),true), 
-					b3BufferInfoCL( convexData.getBufferCL(),true),
-					b3BufferInfoCL( gpuVertices.getBufferCL(),true),
-					b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
-					b3BufferInfoCL( gpuFaces.getBufferCL(),true),
-					b3BufferInfoCL( gpuIndices.getBufferCL(),true),
-					b3BufferInfoCL( clAabbsLocalSpace.getBufferCL(),true),
-					b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
-					b3BufferInfoCL( m_gpuCompoundPairs.getBufferCL()),
-					b3BufferInfoCL( m_numCompoundPairsOut.getBufferCL()),
-					b3BufferInfoCL(subTreesGPU->getBufferCL()),
-					b3BufferInfoCL(treeNodesGPU->getBufferCL()),
-					b3BufferInfoCL(bvhInfo->getBufferCL())
-				};
+		}
+        else
+        {
+            
+           
+            
+            b3AlignedObjectArray<b3Int4> hostPairs;
+            pairs->copyToHost(hostPairs);
+            b3AlignedObjectArray<b3RigidBodyCL> hostBodyBuf;
+            bodyBuf->copyToHost(hostBodyBuf);
 
-				b3LauncherCL launcher(m_queue, m_findCompoundPairsKernel,"m_findCompoundPairsKernel");
-				launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
-				launcher.setConst( nPairs  );
-				launcher.setConst( compoundPairCapacity);
+            b3AlignedObjectArray<b3Collidable> hostCollidables;
+            gpuCollidables.copyToHost(hostCollidables);
+            
+            b3AlignedObjectArray<b3GpuChildShape> cpuChildShapes;
+            gpuChildShapes.copyToHost(cpuChildShapes);
+            
+            b3AlignedObjectArray<b3ConvexPolyhedronCL> hostConvexShapeData;
+            convexData.copyToHost(hostConvexShapeData);
+            
+            b3AlignedObjectArray<b3Vector3> hostVertices;
+            gpuVertices.copyToHost(hostVertices);
+            
+            b3AlignedObjectArray<int> hostHasSepAxis;
+            hostHasSepAxis.resize(nPairs);
+            b3AlignedObjectArray<b3Vector3> hostSepAxis;
+            hostSepAxis.resize(nPairs);
+            
+            b3AlignedObjectArray<b3Vector3> hostUniqueEdges;
+            gpuUniqueEdges.copyToHost(hostUniqueEdges);
+            b3AlignedObjectArray<b3GpuFace> hostFaces;
+            gpuFaces.copyToHost(hostFaces);
+            
+            b3AlignedObjectArray<int> hostIndices;
+            gpuIndices.copyToHost(hostIndices);
+            
+            for (int i=0;i<nPairs;i++)
+            {
+                
+                int bodyIndexA = hostPairs[i].x;
+                int bodyIndexB = hostPairs[i].y;
+                int collidableIndexA = hostBodyBuf[bodyIndexA].m_collidableIdx;
+                int collidableIndexB = hostBodyBuf[bodyIndexB].m_collidableIdx;
+                
+                int shapeIndexA = hostCollidables[collidableIndexA].m_shapeIndex;
+                int shapeIndexB = hostCollidables[collidableIndexB].m_shapeIndex;
+                
+                hostHasSepAxis[i] = 0;
+                
+                //once the broadphase avoids static-static pairs, we can remove this test
+                if ((hostBodyBuf[bodyIndexA].m_invMass==0) &&(hostBodyBuf[bodyIndexB].m_invMass==0))
+                {
+                    continue;
+                }
+                
+                
+                if ((hostCollidables[collidableIndexA].m_shapeType!=SHAPE_CONVEX_HULL) ||(hostCollidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL))
+                {
+                    continue;
+                }
+                
+                float dmin = FLT_MAX;
+                
+                b3ConvexPolyhedronData* convexShapeA = &hostConvexShapeData[shapeIndexA];
+                b3ConvexPolyhedronData* convexShapeB = &hostConvexShapeData[shapeIndexB];
+                b3Vector3 posA = hostBodyBuf[bodyIndexA].m_pos;
+                b3Vector3 posB = hostBodyBuf[bodyIndexB].m_pos;
+                b3Quaternion ornA =hostBodyBuf[bodyIndexA].m_quat;
+                b3Quaternion ornB =hostBodyBuf[bodyIndexB].m_quat;
+                
+                b3Vector3 c0local = hostConvexShapeData[shapeIndexA].m_localCenter;
+                b3Vector3 c0 = b3TransformPoint(c0local, posA, ornA);
+                b3Vector3 c1local = hostConvexShapeData[shapeIndexB].m_localCenter;
+                b3Vector3 c1 = b3TransformPoint(c1local,posB,ornB);
+                b3Vector3 DeltaC2 = c0 - c1;
+                
+                b3Vector3 sepAxis;
+                
+                bool hasSepAxisA = b3FindSeparatingAxis(convexShapeA, convexShapeB, posA, ornA, posB, ornB, DeltaC2,
+                    &hostVertices.at(0), &hostUniqueEdges.at(0), &hostFaces.at(0), &hostIndices.at(0),
+                    &hostVertices.at(0), &hostUniqueEdges.at(0), &hostFaces.at(0), &hostIndices.at(0),
+                                     &sepAxis, &dmin);
+                
+                if (hasSepAxisA)
+                {
+                    bool hasSepAxisB = b3FindSeparatingAxis(convexShapeB, convexShapeA, posB, ornB, posA, ornA, DeltaC2,
+                                                            &hostVertices.at(0), &hostUniqueEdges.at(0), &hostFaces.at(0), &hostIndices.at(0),
+                                                            &hostVertices.at(0), &hostUniqueEdges.at(0), &hostFaces.at(0), &hostIndices.at(0),
+                                                            &sepAxis, &dmin);
+                    if (hasSepAxisB)
+                    {
+                        bool hasEdgeEdge = b3FindSeparatingAxisEdgeEdge(convexShapeA, convexShapeB, posA, ornA, posB, ornB, DeltaC2,
+                                                     &hostVertices.at(0), &hostUniqueEdges.at(0), &hostFaces.at(0), &hostIndices.at(0),
+                                                     &hostVertices.at(0), &hostUniqueEdges.at(0), &hostFaces.at(0), &hostIndices.at(0),
+                                                     &sepAxis, &dmin);
+                        if (hasEdgeEdge)
+                        {
+                            hostHasSepAxis[i] = 1;
+                            hostSepAxis[i] = sepAxis;
+                        }
+                    }
+                }
+            }
+            
+            
+            m_hasSeparatingNormals.copyFromHost(hostHasSepAxis);
+            m_sepNormals.copyFromHost(hostSepAxis);
+        
+            /*
+             //double-check results from GPU (comment-out the 'else' so both paths are executed
+            b3AlignedObjectArray<int> checkHasSepAxis;
+            m_hasSeparatingNormals.copyToHost(checkHasSepAxis);
+            static int frameCount = 0;
+            frameCount++;
+            for (int i=0;i<nPairs;i++)
+            {
+                if (hostHasSepAxis[i] != checkHasSepAxis[i])
+                {
+                    printf("at frameCount %d hostHasSepAxis[%d] = %d but checkHasSepAxis[i] = %d\n",
+                           frameCount,i,hostHasSepAxis[i],checkHasSepAxis[i]);
+                }
+            }
+            //m_hasSeparatingNormals.copyFromHost(hostHasSepAxis);
+            //    m_sepNormals.copyFromHost(hostSepAxis);
+            */
+        }
+        
+        
+        numCompoundPairs = m_numCompoundPairsOut.at(0);
+        bool useGpuFindCompoundPairs=true;
+        if (useGpuFindCompoundPairs)
+        {
+            B3_PROFILE("findCompoundPairsKernel");
+            b3BufferInfoCL bInfo[] = 
+            { 
+                b3BufferInfoCL( pairs->getBufferCL(), true ), 
+                b3BufferInfoCL( bodyBuf->getBufferCL(),true), 
+                b3BufferInfoCL( gpuCollidables.getBufferCL(),true), 
+                b3BufferInfoCL( convexData.getBufferCL(),true),
+                b3BufferInfoCL( gpuVertices.getBufferCL(),true),
+                b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
+                b3BufferInfoCL( gpuFaces.getBufferCL(),true),
+                b3BufferInfoCL( gpuIndices.getBufferCL(),true),
+                b3BufferInfoCL( clAabbsLocalSpace.getBufferCL(),true),
+                b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
+                b3BufferInfoCL( m_gpuCompoundPairs.getBufferCL()),
+                b3BufferInfoCL( m_numCompoundPairsOut.getBufferCL()),
+                b3BufferInfoCL(subTreesGPU->getBufferCL()),
+                b3BufferInfoCL(treeNodesGPU->getBufferCL()),
+                b3BufferInfoCL(bvhInfo->getBufferCL())
+            };
 
-				int num = nPairs;
-				launcher.launch1D( num);
-				clFinish(m_queue);
+            b3LauncherCL launcher(m_queue, m_findCompoundPairsKernel,"m_findCompoundPairsKernel");
+            launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
+            launcher.setConst( nPairs  );
+            launcher.setConst( compoundPairCapacity);
 
-				numCompoundPairs = m_numCompoundPairsOut.at(0);
-				//printf("numCompoundPairs =%d\n",numCompoundPairs );
-				if (numCompoundPairs)
-				{
-					//printf("numCompoundPairs=%d\n",numCompoundPairs);
-				}
-				
+            int num = nPairs;
+            launcher.launch1D( num);
+            clFinish(m_queue);
 
-			} else
-			{
+            numCompoundPairs = m_numCompoundPairsOut.at(0);
+            //printf("numCompoundPairs =%d\n",numCompoundPairs );
+            if (numCompoundPairs)
+            {
+                //printf("numCompoundPairs=%d\n",numCompoundPairs);
+            }
+            
+
+        } else
+        {
 
 
-				b3AlignedObjectArray<b3QuantizedBvhNode>	treeNodesCPU;
-				treeNodesGPU->copyToHost(treeNodesCPU);
+            b3AlignedObjectArray<b3QuantizedBvhNode>	treeNodesCPU;
+            treeNodesGPU->copyToHost(treeNodesCPU);
 
-				b3AlignedObjectArray<b3BvhSubtreeInfo>	subTreesCPU;
-				subTreesGPU->copyToHost(subTreesCPU);
+            b3AlignedObjectArray<b3BvhSubtreeInfo>	subTreesCPU;
+            subTreesGPU->copyToHost(subTreesCPU);
 
-				b3AlignedObjectArray<b3BvhInfo>	bvhInfoCPU;
-				bvhInfo->copyToHost(bvhInfoCPU);
+            b3AlignedObjectArray<b3BvhInfo>	bvhInfoCPU;
+            bvhInfo->copyToHost(bvhInfoCPU);
 
-				b3AlignedObjectArray<b3Aabb> hostAabbsWorldSpace;
-				clAabbsWorldSpace.copyToHost(hostAabbsWorldSpace);
+            b3AlignedObjectArray<b3Aabb> hostAabbsWorldSpace;
+            clAabbsWorldSpace.copyToHost(hostAabbsWorldSpace);
 
-				b3AlignedObjectArray<b3Aabb> hostAabbsLocalSpace;
-				clAabbsLocalSpace.copyToHost(hostAabbsLocalSpace);
+            b3AlignedObjectArray<b3Aabb> hostAabbsLocalSpace;
+            clAabbsLocalSpace.copyToHost(hostAabbsLocalSpace);
 
-				b3AlignedObjectArray<b3Int4> hostPairs;
-				pairs->copyToHost(hostPairs);
+            b3AlignedObjectArray<b3Int4> hostPairs;
+            pairs->copyToHost(hostPairs);
 
-				b3AlignedObjectArray<b3RigidBodyCL> hostBodyBuf;
-				bodyBuf->copyToHost(hostBodyBuf);
+            b3AlignedObjectArray<b3RigidBodyCL> hostBodyBuf;
+            bodyBuf->copyToHost(hostBodyBuf);
 
-				int numCompoundPairsOut=0;
+            int numCompoundPairsOut=0;
 
-				b3AlignedObjectArray<b3Int4> cpuCompoundPairsOut;
-				cpuCompoundPairsOut.resize(compoundPairCapacity);
+            b3AlignedObjectArray<b3Int4> cpuCompoundPairsOut;
+            cpuCompoundPairsOut.resize(compoundPairCapacity);
 
-				b3AlignedObjectArray<b3Collidable> hostCollidables;
-				gpuCollidables.copyToHost(hostCollidables);
-	
-				b3AlignedObjectArray<b3GpuChildShape> cpuChildShapes;
-				gpuChildShapes.copyToHost(cpuChildShapes);
-	
-				b3AlignedObjectArray<b3ConvexPolyhedronCL> hostConvexData;
-				convexData.copyToHost(hostConvexData);
+            b3AlignedObjectArray<b3Collidable> hostCollidables;
+            gpuCollidables.copyToHost(hostCollidables);
 
-				b3AlignedObjectArray<b3Vector3> hostVertices;
-				gpuVertices.copyToHost(hostVertices);
+            b3AlignedObjectArray<b3GpuChildShape> cpuChildShapes;
+            gpuChildShapes.copyToHost(cpuChildShapes);
 
-	
+            b3AlignedObjectArray<b3ConvexPolyhedronCL> hostConvexData;
+            convexData.copyToHost(hostConvexData);
+
+            b3AlignedObjectArray<b3Vector3> hostVertices;
+            gpuVertices.copyToHost(hostVertices);
 
 
-				for (int pairIndex=0;pairIndex<nPairs;pairIndex++)
-				{
-					int bodyIndexA = hostPairs[pairIndex].x;
-					int bodyIndexB = hostPairs[pairIndex].y;
-					int collidableIndexA = hostBodyBuf[bodyIndexA].m_collidableIdx;
-					int collidableIndexB = hostBodyBuf[bodyIndexB].m_collidableIdx;
 
-					findCompoundPairsKernel( 
-								pairIndex,
-								bodyIndexA,
-								bodyIndexB,
-								collidableIndexA,
-								collidableIndexB,
-								&hostBodyBuf[0],
-								&hostCollidables[0],
-								&hostConvexData[0],
-								hostVertices,
-								hostAabbsWorldSpace,
-								hostAabbsLocalSpace,
-								&cpuChildShapes[0],
-								&cpuCompoundPairsOut[0],
-								&numCompoundPairsOut,
-								compoundPairCapacity,
-								treeNodesCPU,
-								subTreesCPU,
-								bvhInfoCPU
-								);
-				}
-				if (numCompoundPairsOut)
-				{
+
+            for (int pairIndex=0;pairIndex<nPairs;pairIndex++)
+            {
+                int bodyIndexA = hostPairs[pairIndex].x;
+                int bodyIndexB = hostPairs[pairIndex].y;
+                int collidableIndexA = hostBodyBuf[bodyIndexA].m_collidableIdx;
+                int collidableIndexB = hostBodyBuf[bodyIndexB].m_collidableIdx;
+
+                findCompoundPairsKernel( 
+                            pairIndex,
+                            bodyIndexA,
+                            bodyIndexB,
+                            collidableIndexA,
+                            collidableIndexB,
+                            &hostBodyBuf[0],
+                            &hostCollidables[0],
+                            &hostConvexData[0],
+                            hostVertices,
+                            hostAabbsWorldSpace,
+                            hostAabbsLocalSpace,
+                            &cpuChildShapes[0],
+                            &cpuCompoundPairsOut[0],
+                            &numCompoundPairsOut,
+                            compoundPairCapacity,
+                            treeNodesCPU,
+                            subTreesCPU,
+                            bvhInfoCPU
+                            );
+            }
+            if (numCompoundPairsOut)
+            {
 //					printf("numCompoundPairsOut=%d\n",numCompoundPairsOut);
-				}
-				
-			}
-			if (numCompoundPairs > compoundPairCapacity)
-			{
-				b3Error("Exceeded compound pair capacity (%d/%d)\n", numCompoundPairs,  compoundPairCapacity);
-				numCompoundPairs = compoundPairCapacity;
-			}
+            }
+            
+        }
+        if (numCompoundPairs > compoundPairCapacity)
+        {
+            b3Error("Exceeded compound pair capacity (%d/%d)\n", numCompoundPairs,  compoundPairCapacity);
+            numCompoundPairs = compoundPairCapacity;
+        }
 
-			
+        
 
-			m_gpuCompoundPairs.resize(numCompoundPairs);
-			m_gpuHasCompoundSepNormals.resize(numCompoundPairs);
-			m_gpuCompoundSepNormals.resize(numCompoundPairs);
-			
+        m_gpuCompoundPairs.resize(numCompoundPairs);
+        m_gpuHasCompoundSepNormals.resize(numCompoundPairs);
+        m_gpuCompoundSepNormals.resize(numCompoundPairs);
+        
 
-			if (numCompoundPairs)
-			{
-				B3_PROFILE("processCompoundPairsPrimitivesKernel");
-				b3BufferInfoCL bInfo[] = 
-				{ 
-					b3BufferInfoCL( m_gpuCompoundPairs.getBufferCL(), true ), 
-					b3BufferInfoCL( bodyBuf->getBufferCL(),true), 
-					b3BufferInfoCL( gpuCollidables.getBufferCL(),true), 
-					b3BufferInfoCL( convexData.getBufferCL(),true),
-					b3BufferInfoCL( gpuVertices.getBufferCL(),true),
-					b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
-					b3BufferInfoCL( gpuFaces.getBufferCL(),true),
-					b3BufferInfoCL( gpuIndices.getBufferCL(),true),
-					b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
-					b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
-					b3BufferInfoCL( contactOut->getBufferCL()),
-					b3BufferInfoCL( m_totalContactsOut.getBufferCL())	
-				};
+        if (numCompoundPairs)
+        {
+            B3_PROFILE("processCompoundPairsPrimitivesKernel");
+            b3BufferInfoCL bInfo[] = 
+            { 
+                b3BufferInfoCL( m_gpuCompoundPairs.getBufferCL(), true ), 
+                b3BufferInfoCL( bodyBuf->getBufferCL(),true), 
+                b3BufferInfoCL( gpuCollidables.getBufferCL(),true), 
+                b3BufferInfoCL( convexData.getBufferCL(),true),
+                b3BufferInfoCL( gpuVertices.getBufferCL(),true),
+                b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
+                b3BufferInfoCL( gpuFaces.getBufferCL(),true),
+                b3BufferInfoCL( gpuIndices.getBufferCL(),true),
+                b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
+                b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
+                b3BufferInfoCL( contactOut->getBufferCL()),
+                b3BufferInfoCL( m_totalContactsOut.getBufferCL())	
+            };
 
-				b3LauncherCL launcher(m_queue, m_processCompoundPairsPrimitivesKernel,"m_processCompoundPairsPrimitivesKernel");
-				launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
-				launcher.setConst( numCompoundPairs  );
-				launcher.setConst(maxContactCapacity);
+            b3LauncherCL launcher(m_queue, m_processCompoundPairsPrimitivesKernel,"m_processCompoundPairsPrimitivesKernel");
+            launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
+            launcher.setConst( numCompoundPairs  );
+            launcher.setConst(maxContactCapacity);
 
-				int num = numCompoundPairs;
-				launcher.launch1D( num);
-				clFinish(m_queue);
-				nContacts = m_totalContactsOut.at(0);
-				//printf("nContacts (after processCompoundPairsPrimitivesKernel) = %d\n",nContacts);
-				if (nContacts>maxContactCapacity)
-				{
-					
-					b3Error("Error: contacts exceeds capacity (%d/%d)\n", nContacts, maxContactCapacity);
-					nContacts = maxContactCapacity;
-				}
-			}
-			
+            int num = numCompoundPairs;
+            launcher.launch1D( num);
+            clFinish(m_queue);
+            nContacts = m_totalContactsOut.at(0);
+            //printf("nContacts (after processCompoundPairsPrimitivesKernel) = %d\n",nContacts);
+            if (nContacts>maxContactCapacity)
+            {
+                
+                b3Error("Error: contacts exceeds capacity (%d/%d)\n", nContacts, maxContactCapacity);
+                nContacts = maxContactCapacity;
+            }
+        }
+        
 
-			if (numCompoundPairs)
-			{
-				B3_PROFILE("processCompoundPairsKernel");
-				b3BufferInfoCL bInfo[] = 
-				{ 
-					b3BufferInfoCL( m_gpuCompoundPairs.getBufferCL(), true ), 
-					b3BufferInfoCL( bodyBuf->getBufferCL(),true), 
-					b3BufferInfoCL( gpuCollidables.getBufferCL(),true), 
-					b3BufferInfoCL( convexData.getBufferCL(),true),
-					b3BufferInfoCL( gpuVertices.getBufferCL(),true),
-					b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
-					b3BufferInfoCL( gpuFaces.getBufferCL(),true),
-					b3BufferInfoCL( gpuIndices.getBufferCL(),true),
-					b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
-					b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
-					b3BufferInfoCL( m_gpuCompoundSepNormals.getBufferCL()),
-					b3BufferInfoCL( m_gpuHasCompoundSepNormals.getBufferCL())
-				};
+        if (numCompoundPairs)
+        {
+            B3_PROFILE("processCompoundPairsKernel");
+            b3BufferInfoCL bInfo[] = 
+            { 
+                b3BufferInfoCL( m_gpuCompoundPairs.getBufferCL(), true ), 
+                b3BufferInfoCL( bodyBuf->getBufferCL(),true), 
+                b3BufferInfoCL( gpuCollidables.getBufferCL(),true), 
+                b3BufferInfoCL( convexData.getBufferCL(),true),
+                b3BufferInfoCL( gpuVertices.getBufferCL(),true),
+                b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
+                b3BufferInfoCL( gpuFaces.getBufferCL(),true),
+                b3BufferInfoCL( gpuIndices.getBufferCL(),true),
+                b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
+                b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
+                b3BufferInfoCL( m_gpuCompoundSepNormals.getBufferCL()),
+                b3BufferInfoCL( m_gpuHasCompoundSepNormals.getBufferCL())
+            };
 
-				b3LauncherCL launcher(m_queue, m_processCompoundPairsKernel,"m_processCompoundPairsKernel");
-				launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
-				launcher.setConst( numCompoundPairs  );
+            b3LauncherCL launcher(m_queue, m_processCompoundPairsKernel,"m_processCompoundPairsKernel");
+            launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
+            launcher.setConst( numCompoundPairs  );
 
-				int num = numCompoundPairs;
-				launcher.launch1D( num);
-				clFinish(m_queue);
-			
-			}
+            int num = numCompoundPairs;
+            launcher.launch1D( num);
+            clFinish(m_queue);
+        
+        }
 
 
-			//printf("numConcave  = %d\n",numConcave);
-
-		}//if (findSeparatingAxisOnGpu)
+        //printf("numConcave  = %d\n",numConcave);
 
+    
 
 //		printf("hostNormals.size()=%d\n",hostNormals.size());
 		//int numPairs = pairCount.at(0);
@@ -3445,8 +3591,8 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
 	
 				clippingFacesOutGPU.resize(numConcavePairs);
 				worldNormalsAGPU.resize(numConcavePairs);
-				worldVertsA1GPU.resize(vertexFaceCapacity*numConcavePairs);
-				worldVertsB1GPU.resize(vertexFaceCapacity*numConcavePairs);
+				worldVertsA1GPU.resize(vertexFaceCapacity*(numConcavePairs));
+				worldVertsB1GPU.resize(vertexFaceCapacity*(numConcavePairs));
 
 
 				if (findConcaveSeparatingAxisKernelGPU)
@@ -3461,34 +3607,118 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
 					*/
 
 					//now perform a SAT test for each triangle-convex element (stored in triangleConvexPairsOut)
-					B3_PROFILE("findConcaveSeparatingAxisKernel");
-					b3BufferInfoCL bInfo[] = { 
-						b3BufferInfoCL( triangleConvexPairsOut.getBufferCL() ), 
-						b3BufferInfoCL( bodyBuf->getBufferCL(),true), 
-						b3BufferInfoCL( gpuCollidables.getBufferCL(),true), 
-						b3BufferInfoCL( convexData.getBufferCL(),true),
-						b3BufferInfoCL( gpuVertices.getBufferCL(),true),
-						b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
-						b3BufferInfoCL( gpuFaces.getBufferCL(),true),
-						b3BufferInfoCL( gpuIndices.getBufferCL(),true),
-						b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
-						b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
-						b3BufferInfoCL( m_concaveSepNormals.getBufferCL()),
-						b3BufferInfoCL( m_concaveHasSeparatingNormals.getBufferCL()),
-						b3BufferInfoCL( clippingFacesOutGPU.getBufferCL()),
-						b3BufferInfoCL( worldVertsA1GPU.getBufferCL()),
-						b3BufferInfoCL(worldNormalsAGPU.getBufferCL()),
-						b3BufferInfoCL(worldVertsB1GPU.getBufferCL())
-					};
+                    if (splitSearchSepAxisConcave)
+                    {
+                        //printf("numConcavePairs = %d\n",numConcavePairs);
+                        m_dmins.resize(numConcavePairs);
+                        {
+                            B3_PROFILE("findConcaveSeparatingAxisVertexFaceKernel");
+                            b3BufferInfoCL bInfo[] = {
+                                b3BufferInfoCL( triangleConvexPairsOut.getBufferCL() ),
+                                b3BufferInfoCL( bodyBuf->getBufferCL(),true),
+                                b3BufferInfoCL( gpuCollidables.getBufferCL(),true),
+                                b3BufferInfoCL( convexData.getBufferCL(),true),
+                                b3BufferInfoCL( gpuVertices.getBufferCL(),true),
+                                b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
+                                b3BufferInfoCL( gpuFaces.getBufferCL(),true),
+                                b3BufferInfoCL( gpuIndices.getBufferCL(),true),
+                                b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
+                                b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
+                                b3BufferInfoCL( m_concaveSepNormals.getBufferCL()),
+                                b3BufferInfoCL( m_concaveHasSeparatingNormals.getBufferCL()),
+                                b3BufferInfoCL( clippingFacesOutGPU.getBufferCL()),
+                                b3BufferInfoCL( worldVertsA1GPU.getBufferCL()),
+                                b3BufferInfoCL(worldNormalsAGPU.getBufferCL()),
+                                b3BufferInfoCL(worldVertsB1GPU.getBufferCL()),
+                                b3BufferInfoCL(m_dmins.getBufferCL())
+                            };
+                            
+                            b3LauncherCL launcher(m_queue, m_findConcaveSeparatingAxisVertexFaceKernel,"m_findConcaveSeparatingAxisVertexFaceKernel");
+                            launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
+                            launcher.setConst(vertexFaceCapacity);
+                            launcher.setConst( numConcavePairs  );
+                            
+                            int num = numConcavePairs;
+                            launcher.launch1D( num);
+                            clFinish(m_queue);
 
-					b3LauncherCL launcher(m_queue, m_findConcaveSeparatingAxisKernel,"m_findConcaveSeparatingAxisKernel");
-					launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
-					launcher.setConst(vertexFaceCapacity);
-					launcher.setConst( numConcavePairs  );
+                            
+                        }
+//                        numConcavePairs = 0;
+                        if (1)
+                        {
+                            B3_PROFILE("findConcaveSeparatingAxisEdgeEdgeKernel");
+                            b3BufferInfoCL bInfo[] = {
+                                b3BufferInfoCL( triangleConvexPairsOut.getBufferCL() ),
+                                b3BufferInfoCL( bodyBuf->getBufferCL(),true),
+                                b3BufferInfoCL( gpuCollidables.getBufferCL(),true),
+                                b3BufferInfoCL( convexData.getBufferCL(),true),
+                                b3BufferInfoCL( gpuVertices.getBufferCL(),true),
+                                b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
+                                b3BufferInfoCL( gpuFaces.getBufferCL(),true),
+                                b3BufferInfoCL( gpuIndices.getBufferCL(),true),
+                                b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
+                                b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
+                                b3BufferInfoCL( m_concaveSepNormals.getBufferCL()),
+                                b3BufferInfoCL( m_concaveHasSeparatingNormals.getBufferCL()),
+                                b3BufferInfoCL( clippingFacesOutGPU.getBufferCL()),
+                                b3BufferInfoCL( worldVertsA1GPU.getBufferCL()),
+                                b3BufferInfoCL(worldNormalsAGPU.getBufferCL()),
+                                b3BufferInfoCL(worldVertsB1GPU.getBufferCL()),
+                                b3BufferInfoCL(m_dmins.getBufferCL())
+                            };
+                            
+                            b3LauncherCL launcher(m_queue, m_findConcaveSeparatingAxisEdgeEdgeKernel,"m_findConcaveSeparatingAxisEdgeEdgeKernel");
+                            launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
+                            launcher.setConst(vertexFaceCapacity);
+                            launcher.setConst( numConcavePairs  );
+                            
+                            int num = numConcavePairs;
+                            launcher.launch1D( num);
+                            clFinish(m_queue);
+                        }
+                      
+                        
+                        // numConcavePairs = 0;
+                        
+                        
+                        
+                        
+                        
+                        
+                    } else
+                    {
+                        B3_PROFILE("findConcaveSeparatingAxisKernel");
+                        b3BufferInfoCL bInfo[] = { 
+                            b3BufferInfoCL( triangleConvexPairsOut.getBufferCL() ), 
+                            b3BufferInfoCL( bodyBuf->getBufferCL(),true), 
+                            b3BufferInfoCL( gpuCollidables.getBufferCL(),true), 
+                            b3BufferInfoCL( convexData.getBufferCL(),true),
+                            b3BufferInfoCL( gpuVertices.getBufferCL(),true),
+                            b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
+                            b3BufferInfoCL( gpuFaces.getBufferCL(),true),
+                            b3BufferInfoCL( gpuIndices.getBufferCL(),true),
+                            b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
+                            b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
+                            b3BufferInfoCL( m_concaveSepNormals.getBufferCL()),
+                            b3BufferInfoCL( m_concaveHasSeparatingNormals.getBufferCL()),
+                            b3BufferInfoCL( clippingFacesOutGPU.getBufferCL()),
+                            b3BufferInfoCL( worldVertsA1GPU.getBufferCL()),
+                            b3BufferInfoCL(worldNormalsAGPU.getBufferCL()),
+                            b3BufferInfoCL(worldVertsB1GPU.getBufferCL())
+                        };
 
-					int num = numConcavePairs;
-					launcher.launch1D( num);
-					clFinish(m_queue);
+                        b3LauncherCL launcher(m_queue, m_findConcaveSeparatingAxisKernel,"m_findConcaveSeparatingAxisKernel");
+                        launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
+                        launcher.setConst(vertexFaceCapacity);
+                        launcher.setConst( numConcavePairs  );
+
+                        int num = numConcavePairs;
+                        launcher.launch1D( num);
+                        clFinish(m_queue);
+                    }
+                    
+                    
 				} else
 				{
 
diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.h b/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.h
index b6f050938..d257d9bc4 100644
--- a/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.h
+++ b/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.h
@@ -31,6 +31,12 @@ struct GpuSatCollision
 	cl_kernel m_findSeparatingAxisEdgeEdgeKernel;
 	
 	cl_kernel				m_findConcaveSeparatingAxisKernel;
+    cl_kernel				m_findConcaveSeparatingAxisVertexFaceKernel;
+    cl_kernel				m_findConcaveSeparatingAxisEdgeEdgeKernel;
+ 
+    
+    
+    
 	cl_kernel				m_findCompoundPairsKernel;
 	cl_kernel				m_processCompoundPairsKernel;
 
diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/sat.cl b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/sat.cl
index 8fea009ee..9a2509425 100644
--- a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/sat.cl
+++ b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/sat.cl
@@ -1381,27 +1381,21 @@ __kernel void   findSeparatingAxisVertexFaceKernel( __global const int4* pairs,
 	
 		int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;
 		int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;
-		
+	
+		hasSeparatingAxis[i] = 0;	
 		
 		//once the broadphase avoids static-static pairs, we can remove this test
 		if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))
 		{
-			hasSeparatingAxis[i] = 0;
 			return;
 		}
 		
 
 		if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONVEX_HULL) ||(collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL))
 		{
-			hasSeparatingAxis[i] = 0;
 			return;
 		}
 			
-		if ((collidables[collidableIndexA].m_shapeType==SHAPE_CONCAVE_TRIMESH))
-		{
-			hasSeparatingAxis[i] = 0;
-			return;
-		}
 
 		int numFacesA = convexShapes[shapeIndexA].m_numFaces;
 
@@ -1524,7 +1518,7 @@ __kernel void   findSeparatingAxisEdgeEdgeKernel( __global const int4* pairs,
 
 
 
-int	findClippingFaces(const float4 separatingNormal,
+inline int	findClippingFaces(const float4 separatingNormal,
                       const ConvexPolyhedronCL* hullA, 
 					  __global const ConvexPolyhedronCL* hullB,
                       const float4 posA, const Quaternion ornA,const float4 posB, const Quaternion ornB,
@@ -1565,11 +1559,17 @@ int	findClippingFaces(const float4 separatingNormal,
     
 	{
 		const btGpuFace polyB = facesB[hullB->m_faceOffset+closestFaceB];
-		const int numVertices = polyB.m_numIndices;
+		int numVertices = polyB.m_numIndices;
+        if (numVertices>capacityWorldVerts)
+            numVertices = capacityWorldVerts;
+        
 		for(int e0=0;e0<numVertices;e0++)
 		{
-			const float4 b = verticesB[hullB->m_vertexOffset+indicesB[polyB.m_indexOffset+e0]];
-			worldVertsB1[pairIndex*capacityWorldVerts+numWorldVertsB1++] = transform(&b,&posB,&ornB);
+            if (e0<capacityWorldVerts)
+            {
+                const float4 b = verticesB[hullB->m_vertexOffset+indicesB[polyB.m_indexOffset+e0]];
+                worldVertsB1[pairIndex*capacityWorldVerts+numWorldVertsB1++] = transform(&b,&posB,&ornB);
+            }
 		}
 	}
     
@@ -1596,10 +1596,16 @@ int	findClippingFaces(const float4 separatingNormal,
 	}
     
     int numVerticesA = facesA[hullA->m_faceOffset+closestFaceA].m_numIndices;
+    if (numVerticesA>capacityWorldVerts)
+       numVerticesA = capacityWorldVerts;
+    
 	for(int e0=0;e0<numVerticesA;e0++)
 	{
-        const float4 a = verticesA[hullA->m_vertexOffset+indicesA[facesA[hullA->m_faceOffset+closestFaceA].m_indexOffset+e0]];
-        worldVertsA1[pairIndex*capacityWorldVerts+e0] = transform(&a, &posA,&ornA);
+        if (e0<capacityWorldVerts)
+        {
+            const float4 a = verticesA[hullA->m_vertexOffset+indicesA[facesA[hullA->m_faceOffset+closestFaceA].m_indexOffset+e0]];
+            worldVertsA1[pairIndex*capacityWorldVerts+e0] = transform(&a, &posA,&ornA);
+        }
     }
     
     clippingFaces[pairIndex].x = closestFaceA;
@@ -1913,3 +1919,543 @@ __kernel void   findConcaveSeparatingAxisKernel( __global int4* concavePairs,
 		concavePairs[pairIdx].w = -1;
 	}
 }
+
+
+
+
+// work-in-progress
+__kernel void   findConcaveSeparatingAxisVertexFaceKernel( __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 const btGpuChildShape* gpuChildShapes,
+                                                __global btAabbCL* aabbs,
+                                                __global float4* concaveSeparatingNormalsOut,
+                                                __global int* concaveHasSeparatingNormals,
+                                                __global int4* clippingFacesOut,
+                                                __global float4* worldVertsA1GPU,
+                                                __global float4*  worldNormalsAGPU,
+                                                __global float4* worldVertsB1GPU,
+                                                __global float* dmins,
+                                                int vertexFaceCapacity,
+                                                int numConcavePairs
+                                                )
+{
+    
+	int i = get_global_id(0);
+	if (i>=numConcavePairs)
+		return;
+    
+	concaveHasSeparatingNormals[i] = 0;
+    
+	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;
+    
+	if (collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL&&
+		collidables[collidableIndexB].m_shapeType!=SHAPE_COMPOUND_OF_CONVEX_HULLS)
+	{
+		concavePairs[pairIdx].w = -1;
+		return;
+	}
+    
+    
+    
+	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 ornA = rigidBodies[bodyIndexA].m_quat;
+		float4 ornB =rigidBodies[bodyIndexB].m_quat;
+        
+		
+        
+        
+		///////////////////
+		///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;
+		}
+		//////////////////
+        
+		float4 c0local = convexPolyhedronA.m_localCenter;
+		float4 c0 = transform(&c0local, &posA, &ornA);
+		float4 c1local = convexShapes[shapeIndexB].m_localCenter;
+		float4 c1 = transform(&c1local,&posB,&ornB);
+		const float4 DeltaC2 = c0 - c1;
+        
+        
+		bool sepA = findSeparatingAxisLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],
+                                             posA,ornA,
+                                             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,
+                                                 posB,ornB,
+                                                 posA,ornA,
+                                                 DeltaC2,
+                                                 vertices,uniqueEdges,faces,indices,
+                                                 verticesA,uniqueEdgesA,facesA,indicesA,
+                                                 &sepAxis,&dmin);
+            
+			if (!sepB)
+			{
+				hasSeparatingAxis = 0;
+			} else
+			{
+				hasSeparatingAxis = 1;
+			}
+		}	
+		
+		if (hasSeparatingAxis)
+		{
+            dmins[i] = dmin;
+			concaveSeparatingNormalsOut[pairIdx]=sepAxis;
+			concaveHasSeparatingNormals[i]=1;
+            
+		} else
+		{	
+			//mark this pair as in-active
+			concavePairs[pairIdx].w = -1;
+		}
+	}
+	else
+	{	
+		//mark this pair as in-active
+		concavePairs[pairIdx].w = -1;
+	}
+}
+
+
+
+
+// work-in-progress
+__kernel void   findConcaveSeparatingAxisEdgeEdgeKernel( __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 const btGpuChildShape* gpuChildShapes,
+                                                          __global btAabbCL* aabbs,
+                                                          __global float4* concaveSeparatingNormalsOut,
+                                                          __global int* concaveHasSeparatingNormals,
+                                                          __global int4* clippingFacesOut,
+                                                          __global float4* worldVertsA1GPU,
+                                                          __global float4*  worldNormalsAGPU,
+                                                          __global float4* worldVertsB1GPU,
+                                                          __global float* dmins,
+                                                          int vertexFaceCapacity,
+                                                          int numConcavePairs
+                                                          )
+{
+    
+	int i = get_global_id(0);
+	if (i>=numConcavePairs)
+		return;
+    
+	if (!concaveHasSeparatingNormals[i])
+        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 = dmins[i];
+		int hasSeparatingAxis=5;
+		float4 sepAxis=make_float4(1,2,3,4);
+        sepAxis = concaveSeparatingNormalsOut[pairIdx];
+        
+		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 ornA = rigidBodies[bodyIndexA].m_quat;
+		float4 ornB =rigidBodies[bodyIndexB].m_quat;
+        
+		
+        
+        
+		///////////////////
+		///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;
+		}
+		//////////////////
+        
+		float4 c0local = convexPolyhedronA.m_localCenter;
+		float4 c0 = transform(&c0local, &posA, &ornA);
+		float4 c1local = convexShapes[shapeIndexB].m_localCenter;
+		float4 c1 = transform(&c1local,&posB,&ornB);
+		const float4 DeltaC2 = c0 - c1;
+        
+        
+		{
+			bool sepEE = findSeparatingAxisEdgeEdgeLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],
+                                                              posA,ornA,
+                                                              posB,ornB,
+                                                              DeltaC2,
+                                                              verticesA,uniqueEdgesA,facesA,indicesA,
+                                                              vertices,uniqueEdges,faces,indices,
+                                                              &sepAxis,&dmin);
+                
+			if (!sepEE)
+			{
+				hasSeparatingAxis = 0;
+			} else
+			{
+				hasSeparatingAxis = 1;
+			}
+		}
+		
+		
+		if (hasSeparatingAxis)
+		{
+			sepAxis.w = dmin;
+            dmins[i] = dmin;
+			concaveSeparatingNormalsOut[pairIdx]=sepAxis;
+			concaveHasSeparatingNormals[i]=1;
+           
+ 	float minDist = -1e30f;
+			float maxDist = 0.02f;
+
+            findClippingFaces(sepAxis,
+                              &convexPolyhedronA,
+                              &convexShapes[shapeIndexB],
+                              posA,ornA,
+                              posB,ornB,
+                              worldVertsA1GPU,
+                              worldNormalsAGPU,
+                              worldVertsB1GPU,
+                              vertexFaceCapacity,
+                              minDist, maxDist,
+                              verticesA,
+                              facesA,
+                              indicesA,
+                              vertices,
+                              faces,
+                              indices,
+                              clippingFacesOut, pairIdx);
+	           
+            
+		} else
+		{	
+			//mark this pair as in-active
+			concavePairs[pairIdx].w = -1;
+		}
+	}
+	else
+	{	
+		//mark this pair as in-active
+		concavePairs[pairIdx].w = -1;
+	}
+}
+
+
+
+
diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satKernels.h b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satKernels.h
index 68df90846..82b2323be 100644
--- a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satKernels.h
+++ b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satKernels.h
@@ -1499,26 +1499,20 @@ static const char* satKernelsCL= \
 "	\n"
 "		int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;\n"
 "		int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;\n"
-"		\n"
+"	\n"
+"		hasSeparatingAxis[i] = 0;	\n"
 "		\n"
 "		//once the broadphase avoids static-static pairs, we can remove this test\n"
 "		if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))\n"
 "		{\n"
-"			hasSeparatingAxis[i] = 0;\n"
 "			return;\n"
 "		}\n"
 "		\n"
 "		if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONVEX_HULL) ||(collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL))\n"
 "		{\n"
-"			hasSeparatingAxis[i] = 0;\n"
 "			return;\n"
 "		}\n"
 "			\n"
-"		if ((collidables[collidableIndexA].m_shapeType==SHAPE_CONCAVE_TRIMESH))\n"
-"		{\n"
-"			hasSeparatingAxis[i] = 0;\n"
-"			return;\n"
-"		}\n"
 "		int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
 "		float dmin = FLT_MAX;\n"
 "		dmins[i] = dmin;\n"
@@ -1627,7 +1621,7 @@ static const char* satKernelsCL= \
 "		}		//if (hasSeparatingAxis[i])\n"
 "	}//(i<numPairs)\n"
 "}\n"
-"int	findClippingFaces(const float4 separatingNormal,\n"
+"inline int	findClippingFaces(const float4 separatingNormal,\n"
 "                      const ConvexPolyhedronCL* hullA, \n"
 "					  __global const ConvexPolyhedronCL* hullB,\n"
 "                      const float4 posA, const Quaternion ornA,const float4 posB, const Quaternion ornB,\n"
@@ -1668,11 +1662,17 @@ static const char* satKernelsCL= \
 "    \n"
 "	{\n"
 "		const btGpuFace polyB = facesB[hullB->m_faceOffset+closestFaceB];\n"
-"		const int numVertices = polyB.m_numIndices;\n"
+"		int numVertices = polyB.m_numIndices;\n"
+"        if (numVertices>capacityWorldVerts)\n"
+"            numVertices = capacityWorldVerts;\n"
+"        \n"
 "		for(int e0=0;e0<numVertices;e0++)\n"
 "		{\n"
-"			const float4 b = verticesB[hullB->m_vertexOffset+indicesB[polyB.m_indexOffset+e0]];\n"
-"			worldVertsB1[pairIndex*capacityWorldVerts+numWorldVertsB1++] = transform(&b,&posB,&ornB);\n"
+"            if (e0<capacityWorldVerts)\n"
+"            {\n"
+"                const float4 b = verticesB[hullB->m_vertexOffset+indicesB[polyB.m_indexOffset+e0]];\n"
+"                worldVertsB1[pairIndex*capacityWorldVerts+numWorldVertsB1++] = transform(&b,&posB,&ornB);\n"
+"            }\n"
 "		}\n"
 "	}\n"
 "    \n"
@@ -1699,10 +1699,16 @@ static const char* satKernelsCL= \
 "	}\n"
 "    \n"
 "    int numVerticesA = facesA[hullA->m_faceOffset+closestFaceA].m_numIndices;\n"
+"    if (numVerticesA>capacityWorldVerts)\n"
+"       numVerticesA = capacityWorldVerts;\n"
+"    \n"
 "	for(int e0=0;e0<numVerticesA;e0++)\n"
 "	{\n"
-"        const float4 a = verticesA[hullA->m_vertexOffset+indicesA[facesA[hullA->m_faceOffset+closestFaceA].m_indexOffset+e0]];\n"
-"        worldVertsA1[pairIndex*capacityWorldVerts+e0] = transform(&a, &posA,&ornA);\n"
+"        if (e0<capacityWorldVerts)\n"
+"        {\n"
+"            const float4 a = verticesA[hullA->m_vertexOffset+indicesA[facesA[hullA->m_faceOffset+closestFaceA].m_indexOffset+e0]];\n"
+"            worldVertsA1[pairIndex*capacityWorldVerts+e0] = transform(&a, &posA,&ornA);\n"
+"        }\n"
 "    }\n"
 "    \n"
 "    clippingFaces[pairIndex].x = closestFaceA;\n"
@@ -1974,4 +1980,531 @@ static const char* satKernelsCL= \
 "		concavePairs[pairIdx].w = -1;\n"
 "	}\n"
 "}\n"
+"// work-in-progress\n"
+"__kernel void   findConcaveSeparatingAxisVertexFaceKernel( __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 const btGpuChildShape* gpuChildShapes,\n"
+"                                                __global btAabbCL* aabbs,\n"
+"                                                __global float4* concaveSeparatingNormalsOut,\n"
+"                                                __global int* concaveHasSeparatingNormals,\n"
+"                                                __global int4* clippingFacesOut,\n"
+"                                                __global float4* worldVertsA1GPU,\n"
+"                                                __global float4*  worldNormalsAGPU,\n"
+"                                                __global float4* worldVertsB1GPU,\n"
+"                                                __global float* dmins,\n"
+"                                                int vertexFaceCapacity,\n"
+"                                                int numConcavePairs\n"
+"                                                )\n"
+"{\n"
+"    \n"
+"	int i = get_global_id(0);\n"
+"	if (i>=numConcavePairs)\n"
+"		return;\n"
+"    \n"
+"	concaveHasSeparatingNormals[i] = 0;\n"
+"    \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"
+"	if (collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL&&\n"
+"		collidables[collidableIndexB].m_shapeType!=SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
+"	{\n"
+"		concavePairs[pairIdx].w = -1;\n"
+"		return;\n"
+"	}\n"
+"    \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"
+"        \n"
+"		float4 ornA = rigidBodies[bodyIndexA].m_quat;\n"
+"		float4 ornB =rigidBodies[bodyIndexB].m_quat;\n"
+"        \n"
+"		\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"
+"		float4 c0local = convexPolyhedronA.m_localCenter;\n"
+"		float4 c0 = transform(&c0local, &posA, &ornA);\n"
+"		float4 c1local = convexShapes[shapeIndexB].m_localCenter;\n"
+"		float4 c1 = transform(&c1local,&posB,&ornB);\n"
+"		const float4 DeltaC2 = c0 - c1;\n"
+"        \n"
+"        \n"
+"		bool sepA = findSeparatingAxisLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],\n"
+"                                             posA,ornA,\n"
+"                                             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,\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"
+"				hasSeparatingAxis = 1;\n"
+"			}\n"
+"		}	\n"
+"		\n"
+"		if (hasSeparatingAxis)\n"
+"		{\n"
+"            dmins[i] = dmin;\n"
+"			concaveSeparatingNormalsOut[pairIdx]=sepAxis;\n"
+"			concaveHasSeparatingNormals[i]=1;\n"
+"            \n"
+"		} else\n"
+"		{	\n"
+"			//mark this pair as in-active\n"
+"			concavePairs[pairIdx].w = -1;\n"
+"		}\n"
+"	}\n"
+"	else\n"
+"	{	\n"
+"		//mark this pair as in-active\n"
+"		concavePairs[pairIdx].w = -1;\n"
+"	}\n"
+"}\n"
+"// work-in-progress\n"
+"__kernel void   findConcaveSeparatingAxisEdgeEdgeKernel( __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 const btGpuChildShape* gpuChildShapes,\n"
+"                                                          __global btAabbCL* aabbs,\n"
+"                                                          __global float4* concaveSeparatingNormalsOut,\n"
+"                                                          __global int* concaveHasSeparatingNormals,\n"
+"                                                          __global int4* clippingFacesOut,\n"
+"                                                          __global float4* worldVertsA1GPU,\n"
+"                                                          __global float4*  worldNormalsAGPU,\n"
+"                                                          __global float4* worldVertsB1GPU,\n"
+"                                                          __global float* dmins,\n"
+"                                                          int vertexFaceCapacity,\n"
+"                                                          int numConcavePairs\n"
+"                                                          )\n"
+"{\n"
+"    \n"
+"	int i = get_global_id(0);\n"
+"	if (i>=numConcavePairs)\n"
+"		return;\n"
+"    \n"
+"	if (!concaveHasSeparatingNormals[i])\n"
+"        return;\n"
+"    \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 = dmins[i];\n"
+"		int hasSeparatingAxis=5;\n"
+"		float4 sepAxis=make_float4(1,2,3,4);\n"
+"        sepAxis = concaveSeparatingNormalsOut[pairIdx];\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"
+"        \n"
+"		float4 ornA = rigidBodies[bodyIndexA].m_quat;\n"
+"		float4 ornB =rigidBodies[bodyIndexB].m_quat;\n"
+"        \n"
+"		\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"
+"		float4 c0local = convexPolyhedronA.m_localCenter;\n"
+"		float4 c0 = transform(&c0local, &posA, &ornA);\n"
+"		float4 c1local = convexShapes[shapeIndexB].m_localCenter;\n"
+"		float4 c1 = transform(&c1local,&posB,&ornB);\n"
+"		const float4 DeltaC2 = c0 - c1;\n"
+"        \n"
+"        \n"
+"		{\n"
+"			bool 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"
+"				hasSeparatingAxis = 0;\n"
+"			} else\n"
+"			{\n"
+"				hasSeparatingAxis = 1;\n"
+"			}\n"
+"		}\n"
+"		\n"
+"		\n"
+"		if (hasSeparatingAxis)\n"
+"		{\n"
+"			sepAxis.w = dmin;\n"
+"            dmins[i] = dmin;\n"
+"			concaveSeparatingNormalsOut[pairIdx]=sepAxis;\n"
+"			concaveHasSeparatingNormals[i]=1;\n"
+"           \n"
+" 	float minDist = -1e30f;\n"
+"			float maxDist = 0.02f;\n"
+"            findClippingFaces(sepAxis,\n"
+"                              &convexPolyhedronA,\n"
+"                              &convexShapes[shapeIndexB],\n"
+"                              posA,ornA,\n"
+"                              posB,ornB,\n"
+"                              worldVertsA1GPU,\n"
+"                              worldNormalsAGPU,\n"
+"                              worldVertsB1GPU,\n"
+"                              vertexFaceCapacity,\n"
+"                              minDist, maxDist,\n"
+"                              verticesA,\n"
+"                              facesA,\n"
+"                              indicesA,\n"
+"                              vertices,\n"
+"                              faces,\n"
+"                              indices,\n"
+"                              clippingFacesOut, pairIdx);\n"
+"	           \n"
+"            \n"
+"		} else\n"
+"		{	\n"
+"			//mark this pair as in-active\n"
+"			concavePairs[pairIdx].w = -1;\n"
+"		}\n"
+"	}\n"
+"	else\n"
+"	{	\n"
+"		//mark this pair as in-active\n"
+"		concavePairs[pairIdx].w = -1;\n"
+"	}\n"
+"}\n"
 ;