diff --git a/Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp b/Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp
index 48d4e3e8a..5fd1f59c0 100644
--- a/Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp
+++ b/Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp
@@ -64,9 +64,15 @@ void	GpuConvexScene::destroyScene()
 int	GpuConvexScene::createDynamicsObjects(const ConstructionInfo& ci)
 {
 	int strideInBytes = 9*sizeof(float);
-	int numVertices = sizeof(barrel_vertices)/strideInBytes;
+	/*int numVertices = sizeof(barrel_vertices)/strideInBytes;
 	int numIndices = sizeof(barrel_indices)/sizeof(int);
 	return createDynamicsObjects2(ci,barrel_vertices,numVertices,barrel_indices,numIndices);
+	*/
+
+	int numVertices = sizeof(tetra_vertices)/strideInBytes;
+	int numIndices = sizeof(tetra_indices)/sizeof(int);
+	return createDynamicsObjects2(ci,tetra_vertices,numVertices,tetra_indices,numIndices);
+	
 }
 
 int	GpuBoxPlaneScene::createDynamicsObjects(const ConstructionInfo& ci)
diff --git a/Demos3/GpuDemos/rigidbody/GpuRigidBodyDemo.cpp b/Demos3/GpuDemos/rigidbody/GpuRigidBodyDemo.cpp
index 95a669cf2..3013c3af2 100644
--- a/Demos3/GpuDemos/rigidbody/GpuRigidBodyDemo.cpp
+++ b/Demos3/GpuDemos/rigidbody/GpuRigidBodyDemo.cpp
@@ -110,7 +110,7 @@ void	GpuRigidBodyDemo::initPhysics(const ConstructionInfo& ci)
 		b3Config config;
 		config.m_maxConvexBodies = b3Max(config.m_maxConvexBodies,ci.arraySizeX*ci.arraySizeY*ci.arraySizeZ+10);
 		config.m_maxConvexShapes = config.m_maxConvexBodies;
-		config.m_maxBroadphasePairs = 12*config.m_maxConvexBodies;
+		config.m_maxBroadphasePairs = 16*config.m_maxConvexBodies;
 		config.m_maxContactCapacity = config.m_maxBroadphasePairs;
 		
 
diff --git a/src/Bullet3Collision/NarrowPhaseCollision/b3Contact4.h b/src/Bullet3Collision/NarrowPhaseCollision/b3Contact4.h
index 6466bae36..13c240468 100644
--- a/src/Bullet3Collision/NarrowPhaseCollision/b3Contact4.h
+++ b/src/Bullet3Collision/NarrowPhaseCollision/b3Contact4.h
@@ -34,6 +34,12 @@ B3_ATTRIBUTE_ALIGNED16(struct) b3Contact4
 	int m_bodyAPtrAndSignBit;
 	int m_bodyBPtrAndSignBit;
 
+	int	m_childIndexA;
+	int	m_childIndexB;
+
+	int m_unused1;
+	int m_unused2;
+
 	int getBodyA()const {return abs(m_bodyAPtrAndSignBit);}
 	int getBodyB()const {return abs(m_bodyBPtrAndSignBit);}
 	bool isBodyAFixed()const { return m_bodyAPtrAndSignBit<0;}
diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.cl b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.cl
index 4eb9339d5..347a2a67a 100644
--- a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.cl
+++ b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.cl
@@ -49,6 +49,12 @@ typedef struct
 
 	int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr
 	int m_bodyBPtrAndSignBit;
+
+	int	m_childIndexA;
+	int	m_childIndexB;
+	int m_unused1;
+	int m_unused2;
+
 } Contact4;
 
 typedef struct 
@@ -483,6 +489,9 @@ void	computeContactSphereConvex(int pairIndex,
 				c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
 				c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
 				c->m_worldPos[0] = pOnB1;
+				c->m_childIndexA = -1;
+				c->m_childIndexB = -1;
+
 				GET_NPOINTS(*c) = 1;
 			} 
 
@@ -705,6 +714,8 @@ void computeContactPlaneConvex(int pairIndex,
 			c->m_batchIdx = pairIndex;
 			c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
 			c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
+			c->m_childIndexA = -1;
+			c->m_childIndexB = -1;
 
 			switch (numReducedPoints)
             {
@@ -783,6 +794,8 @@ void	computeContactPlaneSphere(int pairIndex,
 			c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
 			c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
 			c->m_worldPos[0] = pOnB1;
+			c->m_childIndexA = -1;
+			c->m_childIndexB = -1;
 			GET_NPOINTS(*c) = 1;
 		}//if (dstIdx < numPairs)
 	}//if (hasCollision)
@@ -955,6 +968,8 @@ __kernel void   primitiveContactsKernel( __global const int2* pairs,
 					c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;
 					c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;
 					c->m_worldPos[0] = contactPosB;
+					c->m_childIndexA = -1;
+					c->m_childIndexB = -1;
 					GET_NPOINTS(*c) = 1;
 				}//if (dstIdx < numPairs)
 			}//if ( len <= (radiusA+radiusB))
@@ -1163,7 +1178,8 @@ void	computeContactSphereTriangle(int pairIndex,
 									float4 spherePos2,
 									float radius,
 									float4 pos,
-									float4 quat
+									float4 quat,
+									int faceIndex
 									)
 {
 
@@ -1176,7 +1192,7 @@ void	computeContactSphereTriangle(int pairIndex,
 	float4 closestPnt = (float4)(0, 0, 0, 0);
 	float4 hitNormalWorld = (float4)(0, 0, 0, 0);
 	float minDist = -1000000.f;
-	bool bCollide = true;
+	bool bCollide = false;
 
 	
 	//////////////////////////////////////
@@ -1251,11 +1267,12 @@ void	computeContactSphereTriangle(int pairIndex,
 		closestPnt = contactPoint;
 		float4 contactToCenter = sphereCenter - contactPoint;
 		minDist = length(contactToCenter);
-		if (minDist>0.f)
+		if (minDist>FLT_EPSILON)
 		{
 			hitNormalWorld = normalize(contactToCenter);//*(1./minDist);
+			bCollide  = true;
 		}
-		bCollide  = true;
+		
 	}
 
 
@@ -1273,19 +1290,29 @@ void	computeContactSphereTriangle(int pairIndex,
 		{
 			pOnB1.w = actualDepth;
 			int dstIdx;
-			AppendInc( nGlobalContactsOut, dstIdx );
+
 			
-			if (dstIdx < maxContactCapacity)
+			float lenSqr = dot3F4(normalOnSurfaceB1,normalOnSurfaceB1);
+			if (lenSqr>FLT_EPSILON)
 			{
-				__global Contact4* c = &globalContactsOut[dstIdx];
-				c->m_worldNormal = normalOnSurfaceB1;
-				c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
-				c->m_batchIdx = pairIndex;
-				c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
-				c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
-				c->m_worldPos[0] = pOnB1;
-				GET_NPOINTS(*c) = 1;
-			} 
+				AppendInc( nGlobalContactsOut, dstIdx );
+			
+				if (dstIdx < maxContactCapacity)
+				{
+					__global Contact4* c = &globalContactsOut[dstIdx];
+					c->m_worldNormal = normalOnSurfaceB1;
+					c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
+					c->m_batchIdx = pairIndex;
+					c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
+					c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
+					c->m_worldPos[0] = pOnB1;
+
+					c->m_childIndexA = -1;
+					c->m_childIndexB = faceIndex;
+
+					GET_NPOINTS(*c) = 1;
+				} 
+			}
 
 		}
 	}//if (hasCollision)
@@ -1346,7 +1373,7 @@ __kernel void   findConcaveSphereContactsKernel( __global int4* concavePairs,
 																rigidBodies,collidables,
 																verticesA,
 																globalContactsOut, nGlobalContactsOut,maxContactCapacity,
-																spherePos,sphereRadius,convexPos,convexOrn);
+																spherePos,sphereRadius,convexPos,convexOrn, f);
 
 		return;
 	}
diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.h b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.h
index 84808c5e6..e0f1d1be3 100644
--- a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.h
+++ b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.h
@@ -51,6 +51,12 @@ static const char* primitiveContactsKernelsCL= \
 "\n"
 "	int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr\n"
 "	int m_bodyBPtrAndSignBit;\n"
+"\n"
+"	int	m_childIndexA;\n"
+"	int	m_childIndexB;\n"
+"	int m_unused1;\n"
+"	int m_unused2;\n"
+"\n"
 "} Contact4;\n"
 "\n"
 "typedef struct \n"
@@ -485,6 +491,9 @@ static const char* primitiveContactsKernelsCL= \
 "				c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n"
 "				c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n"
 "				c->m_worldPos[0] = pOnB1;\n"
+"				c->m_childIndexA = -1;\n"
+"				c->m_childIndexB = -1;\n"
+"\n"
 "				GET_NPOINTS(*c) = 1;\n"
 "			} \n"
 "\n"
@@ -707,6 +716,8 @@ static const char* primitiveContactsKernelsCL= \
 "			c->m_batchIdx = pairIndex;\n"
 "			c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n"
 "			c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n"
+"			c->m_childIndexA = -1;\n"
+"			c->m_childIndexB = -1;\n"
 "\n"
 "			switch (numReducedPoints)\n"
 "            {\n"
@@ -785,6 +796,8 @@ static const char* primitiveContactsKernelsCL= \
 "			c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n"
 "			c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n"
 "			c->m_worldPos[0] = pOnB1;\n"
+"			c->m_childIndexA = -1;\n"
+"			c->m_childIndexB = -1;\n"
 "			GET_NPOINTS(*c) = 1;\n"
 "		}//if (dstIdx < numPairs)\n"
 "	}//if (hasCollision)\n"
@@ -957,6 +970,8 @@ static const char* primitiveContactsKernelsCL= \
 "					c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n"
 "					c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n"
 "					c->m_worldPos[0] = contactPosB;\n"
+"					c->m_childIndexA = -1;\n"
+"					c->m_childIndexB = -1;\n"
 "					GET_NPOINTS(*c) = 1;\n"
 "				}//if (dstIdx < numPairs)\n"
 "			}//if ( len <= (radiusA+radiusB))\n"
@@ -1165,7 +1180,8 @@ static const char* primitiveContactsKernelsCL= \
 "									float4 spherePos2,\n"
 "									float radius,\n"
 "									float4 pos,\n"
-"									float4 quat\n"
+"									float4 quat,\n"
+"									int faceIndex\n"
 "									)\n"
 "{\n"
 "\n"
@@ -1178,7 +1194,7 @@ static const char* primitiveContactsKernelsCL= \
 "	float4 closestPnt = (float4)(0, 0, 0, 0);\n"
 "	float4 hitNormalWorld = (float4)(0, 0, 0, 0);\n"
 "	float minDist = -1000000.f;\n"
-"	bool bCollide = true;\n"
+"	bool bCollide = false;\n"
 "\n"
 "	\n"
 "	//////////////////////////////////////\n"
@@ -1253,11 +1269,12 @@ static const char* primitiveContactsKernelsCL= \
 "		closestPnt = contactPoint;\n"
 "		float4 contactToCenter = sphereCenter - contactPoint;\n"
 "		minDist = length(contactToCenter);\n"
-"		if (minDist>0.f)\n"
+"		if (minDist>FLT_EPSILON)\n"
 "		{\n"
 "			hitNormalWorld = normalize(contactToCenter);//*(1./minDist);\n"
+"			bCollide  = true;\n"
 "		}\n"
-"		bCollide  = true;\n"
+"		\n"
 "	}\n"
 "\n"
 "\n"
@@ -1275,19 +1292,29 @@ static const char* primitiveContactsKernelsCL= \
 "		{\n"
 "			pOnB1.w = actualDepth;\n"
 "			int dstIdx;\n"
-"			AppendInc( nGlobalContactsOut, dstIdx );\n"
+"\n"
 "			\n"
-"			if (dstIdx < maxContactCapacity)\n"
+"			float lenSqr = dot3F4(normalOnSurfaceB1,normalOnSurfaceB1);\n"
+"			if (lenSqr>FLT_EPSILON)\n"
 "			{\n"
-"				__global Contact4* c = &globalContactsOut[dstIdx];\n"
-"				c->m_worldNormal = normalOnSurfaceB1;\n"
-"				c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);\n"
-"				c->m_batchIdx = pairIndex;\n"
-"				c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n"
-"				c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n"
-"				c->m_worldPos[0] = pOnB1;\n"
-"				GET_NPOINTS(*c) = 1;\n"
-"			} \n"
+"				AppendInc( nGlobalContactsOut, dstIdx );\n"
+"			\n"
+"				if (dstIdx < maxContactCapacity)\n"
+"				{\n"
+"					__global Contact4* c = &globalContactsOut[dstIdx];\n"
+"					c->m_worldNormal = normalOnSurfaceB1;\n"
+"					c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);\n"
+"					c->m_batchIdx = pairIndex;\n"
+"					c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n"
+"					c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n"
+"					c->m_worldPos[0] = pOnB1;\n"
+"\n"
+"					c->m_childIndexA = -1;\n"
+"					c->m_childIndexB = faceIndex;\n"
+"\n"
+"					GET_NPOINTS(*c) = 1;\n"
+"				} \n"
+"			}\n"
 "\n"
 "		}\n"
 "	}//if (hasCollision)\n"
@@ -1348,7 +1375,7 @@ static const char* primitiveContactsKernelsCL= \
 "																rigidBodies,collidables,\n"
 "																verticesA,\n"
 "																globalContactsOut, nGlobalContactsOut,maxContactCapacity,\n"
-"																spherePos,sphereRadius,convexPos,convexOrn);\n"
+"																spherePos,sphereRadius,convexPos,convexOrn, f);\n"
 "\n"
 "		return;\n"
 "	}\n"
diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.cl b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.cl
index 78f83719f..42cde9b33 100644
--- a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.cl
+++ b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.cl
@@ -50,6 +50,12 @@ typedef struct
 
 	int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr
 	int m_bodyBPtrAndSignBit;
+
+	int	m_childIndexA;
+	int	m_childIndexB;
+	int m_unused1;
+	int m_unused2;
+
 } Contact4;
 
 
@@ -924,6 +930,8 @@ __kernel void   extractManifoldAndAddContactKernel(__global const int2* pairs,
 			int bodyB = pairs[pairIndex].y;
 			c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0 ? -bodyA:bodyA;
 			c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0 ? -bodyB:bodyB;
+			c->m_childIndexA = -1;
+			c->m_childIndexB = -1;
 			for (int i=0;i<nContacts;i++)
 			{
 				c->m_worldPos[i] = localPoints[contactIdx[i]];
@@ -1034,6 +1042,8 @@ __kernel void   clipHullHullKernel( __global const int2* pairs,
 					int bodyB = pairs[pairIndex].y;
 					c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;
 					c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;
+					c->m_childIndexA = -1;
+					c->m_childIndexB = -1;
 
 					for (int i=0;i<nReducedContacts;i++)
 					{
@@ -1165,7 +1175,8 @@ __kernel void   clipCompoundsHullHullKernel( __global const int4* gpuCompoundPai
 					int bodyB = gpuCompoundPairs[pairIndex].y;
 					c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;
 					c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;
-
+					c->m_childIndexA = childShapeIndexA;
+					c->m_childIndexB = childShapeIndexB;
 					for (int i=0;i<nReducedContacts;i++)
 					{
 						c->m_worldPos[i] = pointsIn[contactIdx[i]];
@@ -1241,6 +1252,9 @@ __kernel void   sphereSphereCollisionKernel( __global const int2* pairs,
 					c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;
 					c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;
 					c->m_worldPos[0] = contactPosB;
+					c->m_childIndexA = -1;
+					c->m_childIndexB = -1;
+
 					GET_NPOINTS(*c) = 1;
 				}//if (dstIdx < numPairs)
 			}//if ( len <= (radiusA+radiusB))
@@ -1285,6 +1299,7 @@ __kernel void   clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
 		int bodyIndexA = concavePairsIn[i].x;
 		int bodyIndexB = concavePairsIn[i].y;
 		int f = concavePairsIn[i].z;
+		int childShapeIndexA = f;
 		
 		int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;
 		int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;
@@ -1411,12 +1426,13 @@ __kernel void   clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
 		float4 sepAxis = separatingNormals[i];
 		
 		int shapeTypeB = collidables[collidableIndexB].m_shapeType;
+		int childShapeIndexB =-1;
 		if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS)
 		{
 			///////////////////
 			///compound shape support
 			
-			int childShapeIndexB = concavePairsIn[pairIndex].w;
+			childShapeIndexB = concavePairsIn[pairIndex].w;
 			int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;
 			shapeIndexB = collidables[childColIndexB].m_shapeIndex;
 			float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;
@@ -1468,7 +1484,8 @@ __kernel void   clipHullHullConcaveConvexKernel( __global int4* concavePairsIn,
 				int bodyB = concavePairsIn[pairIndex].y;
 				c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;
 				c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;
-
+				c->m_childIndexA = childShapeIndexA;
+				c->m_childIndexB = childShapeIndexB;
 				for (int i=0;i<nReducedContacts;i++)
 				{
 					c->m_worldPos[i] = pointsIn[contactIdx[i]];
@@ -1888,7 +1905,9 @@ __kernel void   newContactReductionKernel( __global const int2* pairs,
 					int bodyB = pairs[pairIndex].y;
 					c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;
 					c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;
-                    
+                    c->m_childIndexA =-1;
+					c->m_childIndexB =-1;
+
                     switch (nReducedContacts)
                     {
                         case 4:
diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.h b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.h
index 49030b6c7..a4b654925 100644
--- a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.h
+++ b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.h
@@ -52,6 +52,12 @@ static const char* satClipKernelsCL= \
 "\n"
 "	int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr\n"
 "	int m_bodyBPtrAndSignBit;\n"
+"\n"
+"	int	m_childIndexA;\n"
+"	int	m_childIndexB;\n"
+"	int m_unused1;\n"
+"	int m_unused2;\n"
+"\n"
 "} Contact4;\n"
 "\n"
 "\n"
@@ -926,6 +932,8 @@ static const char* satClipKernelsCL= \
 "			int bodyB = pairs[pairIndex].y;\n"
 "			c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0 ? -bodyA:bodyA;\n"
 "			c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0 ? -bodyB:bodyB;\n"
+"			c->m_childIndexA = -1;\n"
+"			c->m_childIndexB = -1;\n"
 "			for (int i=0;i<nContacts;i++)\n"
 "			{\n"
 "				c->m_worldPos[i] = localPoints[contactIdx[i]];\n"
@@ -1036,6 +1044,8 @@ static const char* satClipKernelsCL= \
 "					int bodyB = pairs[pairIndex].y;\n"
 "					c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n"
 "					c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n"
+"					c->m_childIndexA = -1;\n"
+"					c->m_childIndexB = -1;\n"
 "\n"
 "					for (int i=0;i<nReducedContacts;i++)\n"
 "					{\n"
@@ -1167,7 +1177,8 @@ static const char* satClipKernelsCL= \
 "					int bodyB = gpuCompoundPairs[pairIndex].y;\n"
 "					c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n"
 "					c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n"
-"\n"
+"					c->m_childIndexA = childShapeIndexA;\n"
+"					c->m_childIndexB = childShapeIndexB;\n"
 "					for (int i=0;i<nReducedContacts;i++)\n"
 "					{\n"
 "						c->m_worldPos[i] = pointsIn[contactIdx[i]];\n"
@@ -1243,6 +1254,9 @@ static const char* satClipKernelsCL= \
 "					c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n"
 "					c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n"
 "					c->m_worldPos[0] = contactPosB;\n"
+"					c->m_childIndexA = -1;\n"
+"					c->m_childIndexB = -1;\n"
+"\n"
 "					GET_NPOINTS(*c) = 1;\n"
 "				}//if (dstIdx < numPairs)\n"
 "			}//if ( len <= (radiusA+radiusB))\n"
@@ -1287,6 +1301,7 @@ static const char* satClipKernelsCL= \
 "		int bodyIndexA = concavePairsIn[i].x;\n"
 "		int bodyIndexB = concavePairsIn[i].y;\n"
 "		int f = concavePairsIn[i].z;\n"
+"		int childShapeIndexA = f;\n"
 "		\n"
 "		int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
 "		int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
@@ -1413,12 +1428,13 @@ static const char* satClipKernelsCL= \
 "		float4 sepAxis = separatingNormals[i];\n"
 "		\n"
 "		int shapeTypeB = collidables[collidableIndexB].m_shapeType;\n"
+"		int childShapeIndexB =-1;\n"
 "		if (shapeTypeB==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
 "		{\n"
 "			///////////////////\n"
 "			///compound shape support\n"
 "			\n"
-"			int childShapeIndexB = concavePairsIn[pairIndex].w;\n"
+"			childShapeIndexB = concavePairsIn[pairIndex].w;\n"
 "			int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;\n"
 "			shapeIndexB = collidables[childColIndexB].m_shapeIndex;\n"
 "			float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;\n"
@@ -1470,7 +1486,8 @@ static const char* satClipKernelsCL= \
 "				int bodyB = concavePairsIn[pairIndex].y;\n"
 "				c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n"
 "				c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n"
-"\n"
+"				c->m_childIndexA = childShapeIndexA;\n"
+"				c->m_childIndexB = childShapeIndexB;\n"
 "				for (int i=0;i<nReducedContacts;i++)\n"
 "				{\n"
 "					c->m_worldPos[i] = pointsIn[contactIdx[i]];\n"
@@ -1890,7 +1907,9 @@ static const char* satClipKernelsCL= \
 "					int bodyB = pairs[pairIndex].y;\n"
 "					c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n"
 "					c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n"
-"                    \n"
+"                    c->m_childIndexA =-1;\n"
+"					c->m_childIndexB =-1;\n"
+"\n"
 "                    switch (nReducedContacts)\n"
 "                    {\n"
 "                        case 4:\n"
diff --git a/src/Bullet3OpenCL/RigidBody/b3Config.h b/src/Bullet3OpenCL/RigidBody/b3Config.h
index e8b94cfec..be7686ccf 100644
--- a/src/Bullet3OpenCL/RigidBody/b3Config.h
+++ b/src/Bullet3OpenCL/RigidBody/b3Config.h
@@ -29,7 +29,7 @@ struct	b3Config
 		m_maxTriConvexPairCapacity(256*1024)
 	{
 		m_maxConvexShapes = m_maxConvexBodies;
-		m_maxBroadphasePairs = 12*m_maxConvexBodies;
+		m_maxBroadphasePairs = 16*m_maxConvexBodies;
 		m_maxContactCapacity = m_maxBroadphasePairs;
 	}
 };
diff --git a/src/Bullet3OpenCL/RigidBody/b3GpuBatchingPgsSolver.cpp b/src/Bullet3OpenCL/RigidBody/b3GpuBatchingPgsSolver.cpp
index f6eb0ce8e..8b8aa751e 100644
--- a/src/Bullet3OpenCL/RigidBody/b3GpuBatchingPgsSolver.cpp
+++ b/src/Bullet3OpenCL/RigidBody/b3GpuBatchingPgsSolver.cpp
@@ -3,8 +3,9 @@ bool b3GpuBatchContacts = true;
 bool b3GpuSolveConstraint = true;
 bool gpuRadixSort=true;
 bool gpuSetSortData = true;
-bool gpuSortContacts = true;
+
 bool optionalSortContactsDeterminism = true;
+bool gpuSortContactsDeterminism = true;
 
 #include "b3GpuBatchingPgsSolver.h"
 #include "Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.h"
@@ -61,6 +62,10 @@ struct	b3GpuBatchingPgsSolverInternalData
 
 	cl_kernel	m_setDeterminismSortDataBodyAKernel;
 	cl_kernel	m_setDeterminismSortDataBodyBKernel;
+	cl_kernel	m_setDeterminismSortDataChildShapeAKernel;
+	cl_kernel	m_setDeterminismSortDataChildShapeBKernel;
+
+
 
 
 	class b3RadixSort32CL*	m_sort32;
@@ -143,7 +148,9 @@ b3GpuBatchingPgsSolver::b3GpuBatchingPgsSolver(cl_context ctx,cl_device_id devic
 		cl_program solveFrictionProg= b3OpenCLUtils::compileCLProgramFromString( ctx, device, solveFrictionSource, &pErrNum,additionalMacros, B3_SOLVER_FRICTION_KERNEL_PATH);
 		b3Assert(solveFrictionProg);
 
-		cl_program solverSetup2Prog= b3OpenCLUtils::compileCLProgramFromString( ctx, device, solverSetup2Source, &pErrNum,additionalMacros, B3_SOLVER_SETUP2_KERNEL_PATH);
+		//cl_program solverSetup2Prog= b3OpenCLUtils::compileCLProgramFromString( ctx, device, solverSetup2Source, &pErrNum,additionalMacros, B3_SOLVER_SETUP2_KERNEL_PATH);
+		cl_program solverSetup2Prog= b3OpenCLUtils::compileCLProgramFromString( ctx, device, 0, &pErrNum,additionalMacros, B3_SOLVER_SETUP2_KERNEL_PATH,true);
+		
 		b3Assert(solverSetup2Prog);
 
 		
@@ -168,6 +175,13 @@ b3GpuBatchingPgsSolver::b3GpuBatchingPgsSolver(cl_context ctx,cl_device_id devic
 
 		m_data->m_setDeterminismSortDataBodyBKernel =  b3OpenCLUtils::compileCLKernelFromString( ctx, device, solverSetup2Source, "SetDeterminismSortDataBodyB", &pErrNum, solverSetup2Prog,additionalMacros );
 		b3Assert(m_data->m_setDeterminismSortDataBodyBKernel);
+
+		m_data->m_setDeterminismSortDataChildShapeAKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, solverSetup2Source, "SetDeterminismSortDataChildShapeA", &pErrNum, solverSetup2Prog,additionalMacros );
+		b3Assert(m_data->m_setDeterminismSortDataChildShapeAKernel);
+
+		m_data->m_setDeterminismSortDataChildShapeBKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, solverSetup2Source, "SetDeterminismSortDataChildShapeB", &pErrNum, solverSetup2Prog,additionalMacros );
+		b3Assert(m_data->m_setDeterminismSortDataChildShapeBKernel);
+
 		
 		m_data->m_reorderContactKernel = b3OpenCLUtils::compileCLKernelFromString( ctx, device, solverSetup2Source, "ReorderContactKernel", &pErrNum, solverSetup2Prog,additionalMacros );
 		b3Assert(m_data->m_reorderContactKernel);
@@ -234,6 +248,13 @@ b3GpuBatchingPgsSolver::~b3GpuBatchingPgsSolver()
 	clReleaseKernel( m_data->m_reorderContactKernel);
 	clReleaseKernel( m_data->m_copyConstraintKernel);
 
+	clReleaseKernel(m_data->m_setDeterminismSortDataBodyAKernel);
+	clReleaseKernel(m_data->m_setDeterminismSortDataBodyBKernel);
+	clReleaseKernel(m_data->m_setDeterminismSortDataChildShapeAKernel);
+	clReleaseKernel(m_data->m_setDeterminismSortDataChildShapeBKernel);
+
+
+
 	delete m_data;
 }
 
@@ -444,7 +465,11 @@ static bool sortfnc(const b3SortData& a,const b3SortData& b)
 static bool b3ContactCmp(const b3Contact4& p, const b3Contact4& q)
 {
 	return ((p.m_bodyAPtrAndSignBit<q.m_bodyAPtrAndSignBit) ||
-		(p.m_bodyAPtrAndSignBit==q.m_bodyAPtrAndSignBit) && (p.m_bodyBPtrAndSignBit<q.m_bodyBPtrAndSignBit));
+		((p.m_bodyAPtrAndSignBit==q.m_bodyAPtrAndSignBit) && (p.m_bodyBPtrAndSignBit<q.m_bodyBPtrAndSignBit)) ||
+		((p.m_bodyAPtrAndSignBit==q.m_bodyAPtrAndSignBit) && (p.m_bodyBPtrAndSignBit==q.m_bodyBPtrAndSignBit)  &&     p.m_childIndexA<q.m_childIndexA ) ||
+		((p.m_bodyAPtrAndSignBit==q.m_bodyAPtrAndSignBit) && (p.m_bodyBPtrAndSignBit==q.m_bodyBPtrAndSignBit)  &&     p.m_childIndexA<q.m_childIndexA ) ||
+		((p.m_bodyAPtrAndSignBit==q.m_bodyAPtrAndSignBit) && (p.m_bodyBPtrAndSignBit==q.m_bodyBPtrAndSignBit)  &&     p.m_childIndexA==q.m_childIndexA  && p.m_childIndexB<q.m_childIndexB)
+		);
 }
 
 
@@ -552,7 +577,7 @@ void b3GpuBatchingPgsSolver::solveContacts(int numBodies, cl_mem bodyBuf, cl_mem
 
 	if (optionalSortContactsDeterminism)
 	{
-		if (gpuSortContacts)
+		if (gpuSortContactsDeterminism)
 		{
 			B3_PROFILE("GPU Sort contact constraints (determinism)");
 
@@ -561,9 +586,24 @@ void b3GpuBatchingPgsSolver::solveContacts(int numBodies, cl_mem bodyBuf, cl_mem
 
 			m_data->m_pBufContactOutGPU->copyToCL(m_data->m_pBufContactOutGPUCopy->getBufferCL(),numContacts,0,0);
 
-
 			{
-				b3LauncherCL launcher(m_data->m_queue, m_data->m_setDeterminismSortDataBodyAKernel);
+				b3LauncherCL launcher(m_data->m_queue, m_data->m_setDeterminismSortDataChildShapeBKernel);
+				launcher.setBuffer(m_data->m_pBufContactOutGPUCopy->getBufferCL());
+				launcher.setBuffer(m_data->m_contactKeyValues->getBufferCL());
+				launcher.setConst(numContacts);
+				launcher.launch1D( numContacts, 64 );
+			}
+			m_data->m_solverGPU->m_sort32->execute(*m_data->m_contactKeyValues);
+			{
+				b3LauncherCL launcher(m_data->m_queue, m_data->m_setDeterminismSortDataChildShapeAKernel);
+				launcher.setBuffer(m_data->m_pBufContactOutGPUCopy->getBufferCL());
+				launcher.setBuffer(m_data->m_contactKeyValues->getBufferCL());
+				launcher.setConst(numContacts);
+				launcher.launch1D( numContacts, 64 );
+			}
+			m_data->m_solverGPU->m_sort32->execute(*m_data->m_contactKeyValues);
+			{
+				b3LauncherCL launcher(m_data->m_queue, m_data->m_setDeterminismSortDataBodyBKernel);
 				launcher.setBuffer(m_data->m_pBufContactOutGPUCopy->getBufferCL());
 				launcher.setBuffer(m_data->m_contactKeyValues->getBufferCL());
 				launcher.setConst(numContacts);
@@ -573,7 +613,7 @@ void b3GpuBatchingPgsSolver::solveContacts(int numBodies, cl_mem bodyBuf, cl_mem
 			m_data->m_solverGPU->m_sort32->execute(*m_data->m_contactKeyValues);
 			
 			{
-				b3LauncherCL launcher(m_data->m_queue, m_data->m_setDeterminismSortDataBodyBKernel);
+				b3LauncherCL launcher(m_data->m_queue, m_data->m_setDeterminismSortDataBodyAKernel);
 				launcher.setBuffer(m_data->m_pBufContactOutGPUCopy->getBufferCL());
 				launcher.setBuffer(m_data->m_contactKeyValues->getBufferCL());
 				launcher.setConst(numContacts);
@@ -582,8 +622,6 @@ void b3GpuBatchingPgsSolver::solveContacts(int numBodies, cl_mem bodyBuf, cl_mem
 
 			m_data->m_solverGPU->m_sort32->execute(*m_data->m_contactKeyValues);
 
-			//__global Contact4* in, __global Contact4* out, __global int2* sortData, int4 cb )
-
 			{
 				B3_PROFILE("gpu reorderContactKernel (determinism)");
                                 
@@ -886,9 +924,9 @@ void b3GpuBatchingPgsSolver::solveContacts(int numBodies, cl_mem bodyBuf, cl_mem
                                 
 									
 									int simdWidth =numBodies+1;//-1;//64;//-1;//32;
-									int numBatches = sortConstraintByBatch( &cpuContacts[0]+offset, n, simdWidth,csCfg.m_staticIdx ,numBodies);	//	on GPU
+									//int numBatches = sortConstraintByBatch( &cpuContacts[0]+offset, n, simdWidth,csCfg.m_staticIdx ,numBodies);	//	on GPU
 									//int numBatches = sortConstraintByBatch2( &cpuContacts[0]+offset, n, simdWidth,csCfg.m_staticIdx ,numBodies);	//	on GPU
-									//int numBatches = sortConstraintByBatch3( &cpuContacts[0]+offset, n, simdWidth,csCfg.m_staticIdx ,numBodies);	//	on GPU
+									int numBatches = sortConstraintByBatch3( &cpuContacts[0]+offset, n, simdWidth,csCfg.m_staticIdx ,numBodies);	//	on GPU
 									
 									
 
diff --git a/src/Bullet3OpenCL/RigidBody/b3Solver.cpp b/src/Bullet3OpenCL/RigidBody/b3Solver.cpp
index 030187c0a..84537fa19 100644
--- a/src/Bullet3OpenCL/RigidBody/b3Solver.cpp
+++ b/src/Bullet3OpenCL/RigidBody/b3Solver.cpp
@@ -513,6 +513,7 @@ struct SolveTask// : public ThreadPool::Task
 				if (bodyA.m_invMass)
 				{
 					b3Assert(usedBodies[aIdx]==0);
+					usedBodies[aIdx]++;
 				}
 				if (m_wgUsedBodies)
 				{
@@ -537,12 +538,15 @@ struct SolveTask// : public ThreadPool::Task
 						}
 					}
 				}
-				usedBodies[aIdx]++;
+
+
+
 				if (bodyB.m_invMass)
 				{
 					b3Assert(usedBodies[bIdx]==0);
+					usedBodies[bIdx]++;
 				}
-				usedBodies[bIdx]++;
+				
 
 				if( !m_solveFriction )
 				{
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.cl b/src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.cl
index f0efbbe43..ef091021a 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.cl
+++ b/src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.cl
@@ -73,6 +73,12 @@ typedef struct
 
 	int m_bodyA;//sign bit set for fixed objects
 	int m_bodyB;
+
+	int	m_childIndexA;
+	int	m_childIndexB;
+	int m_unused1;
+	int m_unused2;
+
 }Contact4;
 
 typedef struct 
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.h b/src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.h
index 5118951b0..81b5fc667 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.h
+++ b/src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.h
@@ -75,6 +75,12 @@ static const char* batchingKernelsCL= \
 "\n"
 "	int m_bodyA;//sign bit set for fixed objects\n"
 "	int m_bodyB;\n"
+"\n"
+"	int	m_childIndexA;\n"
+"	int	m_childIndexB;\n"
+"	int m_unused1;\n"
+"	int m_unused2;\n"
+"\n"
 "}Contact4;\n"
 "\n"
 "typedef struct \n"
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.cl b/src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.cl
index c6efdf362..fb9361c46 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.cl
+++ b/src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.cl
@@ -74,6 +74,12 @@ typedef struct
 
 	int m_bodyAPtrAndSignBit;//sign bit set for fixed objects
 	int m_bodyBPtrAndSignBit;
+
+	int	m_childIndexA;
+	int	m_childIndexB;
+	int m_unused1;
+	int m_unused2;
+
 }Contact4;
 
 typedef struct 
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.h b/src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.h
index b54f83671..67ff25121 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.h
+++ b/src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.h
@@ -76,6 +76,12 @@ static const char* batchingKernelsNewCL= \
 "\n"
 "	int m_bodyAPtrAndSignBit;//sign bit set for fixed objects\n"
 "	int m_bodyBPtrAndSignBit;\n"
+"\n"
+"	int	m_childIndexA;\n"
+"	int	m_childIndexB;\n"
+"	int m_unused1;\n"
+"	int m_unused2;\n"
+"\n"
 "}Contact4;\n"
 "\n"
 "typedef struct \n"
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.cl b/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.cl
index 1e32c1052..0ad5161bc 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.cl
+++ b/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.cl
@@ -357,12 +357,10 @@ __kernel void breakViolatedConstraintsKernel(__global b3GpuGenericConstraint* co
 	int numRows = numConstraintRows[cid];
 	if (numRows)
 	{
-	//	printf("cid=%d, breakingThreshold =%f\n",cid,breakingThreshold);
 		for (int i=0;i<numRows;i++)
 		{
 			int rowIndex = rowOffsets[cid]+i;
 			float breakingThreshold = constraints[cid].m_breakingImpulseThreshold;
-		//	printf("rows[%d].m_appliedImpulse=%f\n",rowIndex,rows[rowIndex].m_appliedImpulse);
 			if (fabs(rows[rowIndex].m_appliedImpulse) >= breakingThreshold)
 			{
 				constraints[cid].m_flags =0;//&= ~B3_CONSTRAINT_FLAG_ENABLED;
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.h b/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.h
index 83d0fe938..fe08fc8fa 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.h
+++ b/src/Bullet3OpenCL/RigidBody/kernels/jointSolver.h
@@ -359,12 +359,10 @@ static const char* solveConstraintRowsCL= \
 "	int numRows = numConstraintRows[cid];\n"
 "	if (numRows)\n"
 "	{\n"
-"	//	printf(\"cid=%d, breakingThreshold =%f\n\",cid,breakingThreshold);\n"
 "		for (int i=0;i<numRows;i++)\n"
 "		{\n"
 "			int rowIndex = rowOffsets[cid]+i;\n"
 "			float breakingThreshold = constraints[cid].m_breakingImpulseThreshold;\n"
-"		//	printf(\"rows[%d].m_appliedImpulse=%f\n\",rowIndex,rows[rowIndex].m_appliedImpulse);\n"
 "			if (fabs(rows[rowIndex].m_appliedImpulse) >= breakingThreshold)\n"
 "			{\n"
 "				constraints[cid].m_flags =0;//&= ~B3_CONSTRAINT_FLAG_ENABLED;\n"
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/solveContact.cl b/src/Bullet3OpenCL/RigidBody/kernels/solveContact.cl
index 992f20b16..c877dbdd3 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/solveContact.cl
+++ b/src/Bullet3OpenCL/RigidBody/kernels/solveContact.cl
@@ -213,6 +213,12 @@ typedef struct
 
 	int m_bodyAPtrAndSignBit;
 	int m_bodyBPtrAndSignBit;
+	
+	int	m_childIndexA;
+	int	m_childIndexB;
+	int m_unused1;
+	int m_unused2;
+
 } Contact4;
 
 typedef struct
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/solveContact.h b/src/Bullet3OpenCL/RigidBody/kernels/solveContact.h
index 46e375822..f58d8bf76 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/solveContact.h
+++ b/src/Bullet3OpenCL/RigidBody/kernels/solveContact.h
@@ -215,6 +215,12 @@ static const char* solveContactCL= \
 "\n"
 "	int m_bodyAPtrAndSignBit;\n"
 "	int m_bodyBPtrAndSignBit;\n"
+"	\n"
+"	int	m_childIndexA;\n"
+"	int	m_childIndexB;\n"
+"	int m_unused1;\n"
+"	int m_unused2;\n"
+"\n"
 "} Contact4;\n"
 "\n"
 "typedef struct\n"
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/solveFriction.cl b/src/Bullet3OpenCL/RigidBody/kernels/solveFriction.cl
index c0067b0ff..21a302813 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/solveFriction.cl
+++ b/src/Bullet3OpenCL/RigidBody/kernels/solveFriction.cl
@@ -213,6 +213,12 @@ typedef struct
 
 	int m_bodyAPtrAndSignBit;
 	int m_bodyBPtrAndSignBit;
+
+	int	m_childIndexA;
+	int	m_childIndexB;
+	int m_unused1;
+	int m_unused2;
+
 } Contact4;
 
 typedef struct
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/solveFriction.h b/src/Bullet3OpenCL/RigidBody/kernels/solveFriction.h
index 26d41a4d5..01039b960 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/solveFriction.h
+++ b/src/Bullet3OpenCL/RigidBody/kernels/solveFriction.h
@@ -215,6 +215,12 @@ static const char* solveFrictionCL= \
 "\n"
 "	int m_bodyAPtrAndSignBit;\n"
 "	int m_bodyBPtrAndSignBit;\n"
+"\n"
+"	int	m_childIndexA;\n"
+"	int	m_childIndexB;\n"
+"	int m_unused1;\n"
+"	int m_unused2;\n"
+"\n"
 "} Contact4;\n"
 "\n"
 "typedef struct\n"
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/solverSetup.cl b/src/Bullet3OpenCL/RigidBody/kernels/solverSetup.cl
index 814f55646..73e431acd 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/solverSetup.cl
+++ b/src/Bullet3OpenCL/RigidBody/kernels/solverSetup.cl
@@ -412,6 +412,12 @@ typedef struct
 
 	int m_bodyAPtrAndSignBit;
 	int m_bodyBPtrAndSignBit;
+
+	int	m_childIndexA;
+	int	m_childIndexB;
+	int m_unused1;
+	int m_unused2;
+
 } Contact4;
 
 typedef struct
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/solverSetup.h b/src/Bullet3OpenCL/RigidBody/kernels/solverSetup.h
index 83371897b..5c6afc294 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/solverSetup.h
+++ b/src/Bullet3OpenCL/RigidBody/kernels/solverSetup.h
@@ -414,6 +414,12 @@ static const char* solverSetupCL= \
 "\n"
 "	int m_bodyAPtrAndSignBit;\n"
 "	int m_bodyBPtrAndSignBit;\n"
+"\n"
+"	int	m_childIndexA;\n"
+"	int	m_childIndexB;\n"
+"	int m_unused1;\n"
+"	int m_unused2;\n"
+"\n"
 "} Contact4;\n"
 "\n"
 "typedef struct\n"
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.cl b/src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.cl
index 0af8bafaa..5fbc24da8 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.cl
+++ b/src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.cl
@@ -386,6 +386,12 @@ typedef struct
 
 	int m_bodyAPtrAndSignBit;
 	int m_bodyBPtrAndSignBit;
+
+	int	m_childIndexA;
+	int	m_childIndexB;
+	int m_unused1;
+	int m_unused2;
+
 } Contact4;
 
 typedef struct
@@ -441,22 +447,53 @@ void ReorderContactKernel(__global Contact4* in, __global Contact4* out, __globa
 	}
 }
 
-
-__kernel
-__attribute__((reqd_work_group_size(WG_SIZE,1,1)))
-void SetDeterminismSortDataBodyA(__global Contact4* contactsIn, __global int2* sortDataOut, int nContacts)
+__kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1)))
+void SetDeterminismSortDataChildShapeB(__global Contact4* contactsIn, __global int2* sortDataOut, int nContacts)
 {
 	int gIdx = GET_GLOBAL_IDX;
 
 	if( gIdx < nContacts )
 	{
 		int2 sd;
-		sd.x = contactsIn[gIdx].m_bodyAPtrAndSignBit;
+		sd.x = contactsIn[gIdx].m_childIndexB;
 		sd.y = gIdx;
 		sortDataOut[gIdx] = sd;
 	}
 }
 
+__kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1)))
+void SetDeterminismSortDataChildShapeA(__global Contact4* contactsIn, __global int2* sortDataInOut, int nContacts)
+{
+	int gIdx = GET_GLOBAL_IDX;
+
+	if( gIdx < nContacts )
+	{
+		int2 sdIn;
+		sdIn = sortDataInOut[gIdx];
+		int2 sdOut;
+		sdOut.x = contactsIn[sdIn.y].m_childIndexA;
+		sdOut.y = sdIn.y;
+		sortDataInOut[gIdx] = sdOut;
+	}
+}
+
+__kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1)))
+void SetDeterminismSortDataBodyA(__global Contact4* contactsIn, __global int2* sortDataInOut, int nContacts)
+{
+	int gIdx = GET_GLOBAL_IDX;
+
+	if( gIdx < nContacts )
+	{
+		int2 sdIn;
+		sdIn = sortDataInOut[gIdx];
+		int2 sdOut;
+		sdOut.x = contactsIn[sdIn.y].m_bodyAPtrAndSignBit;
+		sdOut.y = sdIn.y;
+		sortDataInOut[gIdx] = sdOut;
+	}
+}
+
+
 __kernel
 __attribute__((reqd_work_group_size(WG_SIZE,1,1)))
 void SetDeterminismSortDataBodyB(__global Contact4* contactsIn, __global int2* sortDataInOut, int nContacts)
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.h b/src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.h
index 39b1e158b..d7ba94a2d 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.h
+++ b/src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.h
@@ -388,6 +388,12 @@ static const char* solverSetup2CL= \
 "\n"
 "	int m_bodyAPtrAndSignBit;\n"
 "	int m_bodyBPtrAndSignBit;\n"
+"\n"
+"	int	m_childIndexA;\n"
+"	int	m_childIndexB;\n"
+"	int m_unused1;\n"
+"	int m_unused2;\n"
+"\n"
 "} Contact4;\n"
 "\n"
 "typedef struct\n"
@@ -443,22 +449,53 @@ static const char* solverSetup2CL= \
 "	}\n"
 "}\n"
 "\n"
-"\n"
-"__kernel\n"
-"__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
-"void SetDeterminismSortDataBodyA(__global Contact4* contactsIn, __global int2* sortDataOut, int nContacts)\n"
+"__kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
+"void SetDeterminismSortDataChildShapeA(__global Contact4* contactsIn, __global int2* sortDataOut, int nContacts)\n"
 "{\n"
 "	int gIdx = GET_GLOBAL_IDX;\n"
 "\n"
 "	if( gIdx < nContacts )\n"
 "	{\n"
 "		int2 sd;\n"
-"		sd.x = contactsIn[gIdx].m_bodyAPtrAndSignBit;\n"
+"		sd.x = contactsIn[gIdx].m_childIndexA;\n"
 "		sd.y = gIdx;\n"
 "		sortDataOut[gIdx] = sd;\n"
 "	}\n"
 "}\n"
 "\n"
+"__kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
+"void SetDeterminismSortDataChildShapeB(__global Contact4* contactsIn, __global int2* sortDataInOut, int nContacts)\n"
+"{\n"
+"	int gIdx = GET_GLOBAL_IDX;\n"
+"\n"
+"	if( gIdx < nContacts )\n"
+"	{\n"
+"		int2 sdIn;\n"
+"		sdIn = sortDataInOut[gIdx];\n"
+"		int2 sdOut;\n"
+"		sdOut.x = contactsIn[sdIn.y].m_childIndexB;\n"
+"		sdOut.y = sdIn.y;\n"
+"		sortDataInOut[gIdx] = sdOut;\n"
+"	}\n"
+"}\n"
+"\n"
+"__kernel __attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
+"void SetDeterminismSortDataBodyA(__global Contact4* contactsIn, __global int2* sortDataInOut, int nContacts)\n"
+"{\n"
+"	int gIdx = GET_GLOBAL_IDX;\n"
+"\n"
+"	if( gIdx < nContacts )\n"
+"	{\n"
+"		int2 sdIn;\n"
+"		sdIn = sortDataInOut[gIdx];\n"
+"		int2 sdOut;\n"
+"		sdOut.x = contactsIn[sdIn.y].m_bodyAPtrAndSignBit;\n"
+"		sdOut.y = sdIn.y;\n"
+"		sortDataInOut[gIdx] = sdOut;\n"
+"	}\n"
+"}\n"
+"\n"
+"\n"
 "__kernel\n"
 "__attribute__((reqd_work_group_size(WG_SIZE,1,1)))\n"
 "void SetDeterminismSortDataBodyB(__global Contact4* contactsIn, __global int2* sortDataInOut, int nContacts)\n"
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.cl b/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.cl
index 933f39e88..19e58220f 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.cl
+++ b/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.cl
@@ -389,6 +389,12 @@ typedef struct
 
 	int m_bodyAPtrAndSignBit;
 	int m_bodyBPtrAndSignBit;
+
+	int	m_childIndexA;
+	int	m_childIndexB;
+	int m_unused1;
+	int m_unused2;
+
 } Contact4;
 
 
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h b/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h
index 29cd712a1..2c31dc31d 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h
+++ b/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h
@@ -391,6 +391,12 @@ static const char* solverUtilsCL= \
 "\n"
 "	int m_bodyAPtrAndSignBit;\n"
 "	int m_bodyBPtrAndSignBit;\n"
+"\n"
+"	int	m_childIndexA;\n"
+"	int	m_childIndexB;\n"
+"	int m_unused1;\n"
+"	int m_unused2;\n"
+"\n"
 "} Contact4;\n"
 "\n"
 "\n"