diff --git a/src/Bullet3Collision/BroadPhaseCollision/b3OverlappingPair.h b/src/Bullet3Collision/BroadPhaseCollision/b3OverlappingPair.h
index a27cb620c..611e183bc 100644
--- a/src/Bullet3Collision/BroadPhaseCollision/b3OverlappingPair.h
+++ b/src/Bullet3Collision/BroadPhaseCollision/b3OverlappingPair.h
@@ -16,7 +16,7 @@ subject to the following restrictions:
 #ifndef B3_OVERLAPPING_PAIR_H
 #define B3_OVERLAPPING_PAIR_H
 
-#include "Bullet3Common/b3Int4.h"
+#include "Bullet3Common/shared/b3Int4.h"
 
 #define B3_NEW_PAIR_MARKER -1
 #define B3_REMOVED_PAIR_MARKER -2
diff --git a/src/Bullet3Collision/BroadPhaseCollision/b3OverlappingPairCache.h b/src/Bullet3Collision/BroadPhaseCollision/b3OverlappingPairCache.h
index 35099be37..ae0799fbf 100644
--- a/src/Bullet3Collision/BroadPhaseCollision/b3OverlappingPairCache.h
+++ b/src/Bullet3Collision/BroadPhaseCollision/b3OverlappingPairCache.h
@@ -16,7 +16,7 @@ subject to the following restrictions:
 #ifndef B3_OVERLAPPING_PAIR_CACHE_H
 #define B3_OVERLAPPING_PAIR_CACHE_H
 
-#include "Bullet3Common/b3Int2.h"
+#include "Bullet3Common/shared/b3Int2.h"
 #include "Bullet3Common/b3AlignedObjectArray.h"
 
 class b3Dispatcher;
diff --git a/src/Bullet3Collision/NarrowPhaseCollision/b3RigidBodyCL.h b/src/Bullet3Collision/NarrowPhaseCollision/b3RigidBodyCL.h
index 9401eb28a..28d5ab4f4 100644
--- a/src/Bullet3Collision/NarrowPhaseCollision/b3RigidBodyCL.h
+++ b/src/Bullet3Collision/NarrowPhaseCollision/b3RigidBodyCL.h
@@ -18,20 +18,14 @@ subject to the following restrictions:
 
 #include "Bullet3Common/b3Scalar.h"
 #include "Bullet3Common/b3Matrix3x3.h"
+#include "Bullet3Collision/NarrowPhaseCollision/Shared/b3RigidBodyData.h"
 
-B3_ATTRIBUTE_ALIGNED16(struct) b3RigidBodyCL
+
+B3_ATTRIBUTE_ALIGNED16(struct) b3RigidBodyCL : public b3RigidBodyData
 {
 	B3_DECLARE_ALIGNED_ALLOCATOR();
 
-	b3Vector3 		m_pos;
-	b3Quaternion 	m_quat;
-	b3Vector3			m_linVel;
-	b3Vector3			m_angVel;
-
-	int 					m_collidableIdx;
-	float 				m_invMass;
-	float 				m_restituitionCoeff;
-	float 				m_frictionCoeff;
+	
 
 	float	getInvMass() const
 	{
diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/b3Collidable.h b/src/Bullet3Collision/NarrowPhaseCollision/shared/b3Collidable.h
similarity index 71%
rename from src/Bullet3OpenCL/NarrowphaseCollision/b3Collidable.h
rename to src/Bullet3Collision/NarrowPhaseCollision/shared/b3Collidable.h
index a1c923655..c3407a769 100644
--- a/src/Bullet3OpenCL/NarrowphaseCollision/b3Collidable.h
+++ b/src/Bullet3Collision/NarrowPhaseCollision/shared/b3Collidable.h
@@ -2,8 +2,9 @@
 #ifndef B3_COLLIDABLE_H
 #define B3_COLLIDABLE_H
 
-#include "Bullet3Common/b3Vector3.h"
-#include "Bullet3Common/b3Quaternion.h"
+
+#include "Bullet3Common/shared/b3Float4.h"
+#include "Bullet3Common/shared/b3Quat.h"
 
 enum b3ShapeTypes
 {
@@ -23,22 +24,22 @@ struct b3Collidable
 		int m_numChildShapes;
 		int m_bvhIndex;
 	};
-	float m_radius;
+	union
+	{
+		float m_radius;
+		int	m_compoundBvhIndex;
+	};
+
 	int m_shapeType;
 	int m_shapeIndex;
 };
 
-struct b3CollidableNew
-{
-	short int m_shapeType;
-	short int m_numShapes;
-	int m_shapeIndex;
-};
+
 
 struct b3GpuChildShape
 {
-	b3Vector3	m_childPosition;
-	b3Quaternion m_childOrientation;
+	b3Float4	m_childPosition;
+	b3Quat		m_childOrientation;
 	int m_shapeIndex;
 	int m_unused0;
 	int m_unused1;
diff --git a/src/Bullet3Collision/NarrowPhaseCollision/shared/b3CollidableData.h b/src/Bullet3Collision/NarrowPhaseCollision/shared/b3CollidableData.h
new file mode 100644
index 000000000..e69de29bb
diff --git a/src/Bullet3Collision/NarrowPhaseCollision/shared/b3ConvexPolyhedronData.h b/src/Bullet3Collision/NarrowPhaseCollision/shared/b3ConvexPolyhedronData.h
new file mode 100644
index 000000000..d0aceb6c1
--- /dev/null
+++ b/src/Bullet3Collision/NarrowPhaseCollision/shared/b3ConvexPolyhedronData.h
@@ -0,0 +1,37 @@
+
+#ifndef B3_CONVEX_POLYHEDRON_DATA_H
+#define B3_CONVEX_POLYHEDRON_DATA_H
+
+
+
+#include "Bullet3Common/shared/b3Float4.h"
+#include "Bullet3Common/shared/b3Quat.h"
+
+struct b3GpuFace
+{
+	b3Float4 m_plane;
+	int m_indexOffset;
+	int m_numIndices;
+	int m_unusedPadding1;
+	int m_unusedPadding2;
+};
+
+B3_ATTRIBUTE_ALIGNED16(struct) b3ConvexPolyhedronData
+{
+	b3Float4		m_localCenter;
+	b3Float4		m_extents;
+	b3Float4		mC;
+	b3Float4		mE;
+
+	float			m_radius;
+	int	m_faceOffset;
+	int m_numFaces;
+	int	m_numVertices;
+
+	int m_vertexOffset;
+	int	m_uniqueEdgesOffset;
+	int	m_numUniqueEdges;
+	int m_unused;
+};
+
+#endif //B3_CONVEX_POLYHEDRON_DATA_H
diff --git a/src/Bullet3Collision/NarrowPhaseCollision/shared/b3RigidBodyData.h b/src/Bullet3Collision/NarrowPhaseCollision/shared/b3RigidBodyData.h
new file mode 100644
index 000000000..473e5cc44
--- /dev/null
+++ b/src/Bullet3Collision/NarrowPhaseCollision/shared/b3RigidBodyData.h
@@ -0,0 +1,21 @@
+#ifndef B3_RIGIDBODY_DATA_H
+#define B3_RIGIDBODY_DATA_H
+
+#include "Bullet3Common/shared/b3Float4.h"
+#include "Bullet3Common/shared/b3Quat.h"
+
+struct b3RigidBodyData
+{
+	b3Float4				m_pos;
+	b3Quat					m_quat;
+	b3Float4				m_linVel;
+	b3Float4				m_angVel;
+
+	int 					m_collidableIdx;
+	float 				m_invMass;
+	float 				m_restituitionCoeff;
+	float 				m_frictionCoeff;
+};
+	
+#endif //B3_RIGIDBODY_DATA_H
+	
\ No newline at end of file
diff --git a/src/Bullet3Common/b3Quaternion.h b/src/Bullet3Common/b3Quaternion.h
index 1472dbe6f..bde9288bd 100644
--- a/src/Bullet3Common/b3Quaternion.h
+++ b/src/Bullet3Common/b3Quaternion.h
@@ -832,6 +832,12 @@ b3Slerp(const b3Quaternion& q1, const b3Quaternion& q2, const b3Scalar& t)
 	return q1.slerp(q2, t);
 }
 
+B3_FORCE_INLINE b3Quaternion
+b3QuatMul(const b3Quaternion& rot0, const b3Quaternion& rot1)
+{
+	return rot0*rot1;
+}
+
 B3_FORCE_INLINE b3Vector3 
 b3QuatRotate(const b3Quaternion& rotation, const b3Vector3& v) 
 {
diff --git a/src/Bullet3Common/shared/b3Float4.h b/src/Bullet3Common/shared/b3Float4.h
index b5006a8fe..ceb7c1e5a 100644
--- a/src/Bullet3Common/shared/b3Float4.h
+++ b/src/Bullet3Common/shared/b3Float4.h
@@ -4,8 +4,9 @@
 #include "Bullet3Common/shared/b3PlatformDefinitions.h"
 
 #ifdef __cplusplus
+	#include "Bullet3Common/b3Vector3.h"
 	#define b3Float4 b3Vector3
-#else//bla
+#else
 	typedef float4	b3Float4;
 #endif 
 
diff --git a/src/Bullet3Common/b3Int2.h b/src/Bullet3Common/shared/b3Int2.h
similarity index 100%
rename from src/Bullet3Common/b3Int2.h
rename to src/Bullet3Common/shared/b3Int2.h
diff --git a/src/Bullet3Common/b3Int4.h b/src/Bullet3Common/shared/b3Int4.h
similarity index 100%
rename from src/Bullet3Common/b3Int4.h
rename to src/Bullet3Common/shared/b3Int4.h
diff --git a/src/Bullet3Common/shared/b3PlatformDefinitions.h b/src/Bullet3Common/shared/b3PlatformDefinitions.h
index 3dccd8431..64f784e8a 100644
--- a/src/Bullet3Common/shared/b3PlatformDefinitions.h
+++ b/src/Bullet3Common/shared/b3PlatformDefinitions.h
@@ -6,4 +6,11 @@ struct MyTest
 	int bla;
 };
 
+#ifdef __cplusplus
+#define b3AtomicInc(a) ((*a)++)
+#else
+#define b3AtomicInc atomic_inc
+
+#endif
+
 #endif
diff --git a/src/Bullet3Common/shared/b3Quat.h b/src/Bullet3Common/shared/b3Quat.h
new file mode 100644
index 000000000..3bfe40a46
--- /dev/null
+++ b/src/Bullet3Common/shared/b3Quat.h
@@ -0,0 +1,13 @@
+#ifndef B3_QUAT_H
+#define B3_QUAT_H
+
+#include "Bullet3Common/shared/b3PlatformDefinitions.h"
+
+#ifdef __cplusplus
+	#include "Bullet3Common/b3Quaternion.h"
+	#define b3Quat b3Quaternion
+#else
+	typedef float4	b3Quat;
+#endif 
+
+#endif //B3_QUAT_H
diff --git a/src/Bullet3Dynamics/ConstraintSolver/b3SolverBody.h b/src/Bullet3Dynamics/ConstraintSolver/b3SolverBody.h
index 8ec5ec7aa..2cb23bb49 100644
--- a/src/Bullet3Dynamics/ConstraintSolver/b3SolverBody.h
+++ b/src/Bullet3Dynamics/ConstraintSolver/b3SolverBody.h
@@ -16,7 +16,7 @@ subject to the following restrictions:
 #ifndef B3_SOLVER_BODY_H
 #define B3_SOLVER_BODY_H
 
-class	b3RigidBody;
+struct	b3RigidBody;
 #include "Bullet3Common/b3Vector3.h"
 #include "Bullet3Common/b3Matrix3x3.h"
 
diff --git a/src/Bullet3Dynamics/ConstraintSolver/b3SolverConstraint.h b/src/Bullet3Dynamics/ConstraintSolver/b3SolverConstraint.h
index ad7e2da52..edeadb459 100644
--- a/src/Bullet3Dynamics/ConstraintSolver/b3SolverConstraint.h
+++ b/src/Bullet3Dynamics/ConstraintSolver/b3SolverConstraint.h
@@ -16,7 +16,7 @@ subject to the following restrictions:
 #ifndef B3_SOLVER_CONSTRAINT_H
 #define B3_SOLVER_CONSTRAINT_H
 
-class	b3RigidBody;
+struct b3RigidBody;
 #include "Bullet3Common/b3Vector3.h"
 #include "Bullet3Common/b3Matrix3x3.h"
 //#include "b3JacobianEntry.h"
diff --git a/src/Bullet3Dynamics/ConstraintSolver/b3TypedConstraint.h b/src/Bullet3Dynamics/ConstraintSolver/b3TypedConstraint.h
index 0bded5b0c..8069a556d 100644
--- a/src/Bullet3Dynamics/ConstraintSolver/b3TypedConstraint.h
+++ b/src/Bullet3Dynamics/ConstraintSolver/b3TypedConstraint.h
@@ -61,6 +61,7 @@ B3_ATTRIBUTE_ALIGNED16(struct)	b3JointFeedback
 	b3Vector3	m_appliedTorqueBodyB;
 };
 
+
 struct b3RigidBodyCL;
 
 
diff --git a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h
index be6a8740a..7d113f23c 100644
--- a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h
+++ b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h
@@ -7,7 +7,7 @@ class b3Vector3;
 #include "Bullet3OpenCL/ParallelPrimitives/b3RadixSort32CL.h"
 
 #include "b3SapAabb.h"
-#include "Bullet3Common/b3Int2.h"
+#include "Bullet3Common/shared/b3Int2.h"
 
 
 class b3GpuSapBroadphase
diff --git a/src/Bullet3OpenCL/BroadphaseCollision/b3SapAabb.h b/src/Bullet3OpenCL/BroadphaseCollision/b3SapAabb.h
index 7f971ba54..ea6550fed 100644
--- a/src/Bullet3OpenCL/BroadphaseCollision/b3SapAabb.h
+++ b/src/Bullet3OpenCL/BroadphaseCollision/b3SapAabb.h
@@ -2,19 +2,13 @@
 #define B3_SAP_AABB_H
 
 #include "Bullet3Common/b3Scalar.h"
+#include "Bullet3Collision/BroadPhaseCollision/shared/b3Aabb.h"
 
-B3_ATTRIBUTE_ALIGNED16(struct) b3SapAabb
+///just make sure that the b3Aabb is 16-byte aligned
+B3_ATTRIBUTE_ALIGNED16(struct) b3SapAabb : public b3Aabb
 {
-	union
-	{
-		float m_min[4];
-		int m_minIndices[4];
-	};
-	union
-	{
-		float m_max[4];
-		int m_signedMaxIndices[4];
-	};
+
 };
 
+
 #endif //B3_SAP_AABB_H
diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.cpp b/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.cpp
index 6d1e43b5e..ef66cf34b 100644
--- a/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.cpp
+++ b/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.cpp
@@ -31,6 +31,7 @@ int b3g_actualSATPairTests=0;
 #include <string.h>//memcpy
 #include "b3ConvexPolyhedronCL.h"
 #include "Bullet3OpenCL/NarrowphaseCollision/b3ContactCache.h"
+#include "Bullet3Geometry/b3AabbUtil.h"
 
 typedef b3AlignedObjectArray<b3Vector3> b3VertexArray;
 
@@ -54,6 +55,15 @@ typedef b3AlignedObjectArray<b3Vector3> b3VertexArray;
 #define BT_NARROWPHASE_PRIMITIVE_CONTACT_PATH "src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.cl"
 
 
+#ifndef __global
+#define __global
+#endif
+
+#ifndef __kernel
+#define __kernel
+#endif
+
+
 #define dot3F4 b3Dot
 
 GpuSatCollision::GpuSatCollision(cl_context ctx,cl_device_id device, cl_command_queue  q )
@@ -419,8 +429,809 @@ int extractManifoldSequentialGlobal( const float4* p, int nPoints, const float4&
 
 
 
+#define MAX_VERTS 1024
 
 
+inline void project(const b3ConvexPolyhedronData& hull,  const float4& pos, const b3Quaternion& orn, const float4& dir, const b3AlignedObjectArray<b3Vector3>& vertices, b3Scalar& min, b3Scalar& max)
+{
+	min = FLT_MAX;
+	max = -FLT_MAX;
+	int numVerts = hull.m_numVertices;
+
+	const float4 localDir = b3QuatRotate(orn.inverse(),dir);
+
+	b3Scalar offset = dot3F4(pos,dir);
+
+	for(int i=0;i<numVerts;i++)
+	{
+		//b3Vector3 pt = trans * vertices[m_vertexOffset+i];
+		//b3Scalar dp = pt.dot(dir);
+		b3Vector3 vertex = vertices[hull.m_vertexOffset+i];
+		b3Scalar dp = dot3F4((float4&)vertices[hull.m_vertexOffset+i],localDir);
+		//b3Assert(dp==dpL);
+		if(dp < min)	min = dp;
+		if(dp > max)	max = dp;
+	}
+	if(min>max)
+	{
+		b3Scalar tmp = min;
+		min = max;
+		max = tmp;
+	}
+	min += offset;
+	max += offset;
+}
+
+
+static bool TestSepAxis(const b3ConvexPolyhedronData& hullA, const b3ConvexPolyhedronData& hullB, 
+	const float4& posA,const b3Quaternion& ornA,
+	const float4& posB,const b3Quaternion& ornB,
+	const float4& sep_axis, const b3AlignedObjectArray<b3Vector3>& verticesA,const b3AlignedObjectArray<b3Vector3>& verticesB,b3Scalar& depth)
+{
+	b3Scalar Min0,Max0;
+	b3Scalar Min1,Max1;
+	project(hullA,posA,ornA,sep_axis,verticesA, Min0, Max0);
+	project(hullB,posB,ornB, sep_axis,verticesB, Min1, Max1);
+
+	if(Max0<Min1 || Max1<Min0)
+		return false;
+
+	b3Scalar d0 = Max0 - Min1;
+	assert(d0>=0.0f);
+	b3Scalar d1 = Max1 - Min0;
+	assert(d1>=0.0f);
+	depth = d0<d1 ? d0:d1;
+	return true;
+}
+
+inline bool IsAlmostZero(const b3Vector3& v)
+{
+	if(fabsf(v.x)>1e-6 || fabsf(v.y)>1e-6 || fabsf(v.z)>1e-6)	return false;
+	return true;
+}
+
+
+static bool findSeparatingAxis(	const b3ConvexPolyhedronData& hullA, const b3ConvexPolyhedronData& hullB, 
+	const float4& posA1,
+	const b3Quaternion& ornA,
+	const float4& posB1,
+	const b3Quaternion& ornB,
+	const b3AlignedObjectArray<b3Vector3>& verticesA,
+	const b3AlignedObjectArray<b3Vector3>& uniqueEdgesA, 
+	const b3AlignedObjectArray<b3GpuFace>& facesA,
+	const b3AlignedObjectArray<int>& indicesA,
+	const b3AlignedObjectArray<b3Vector3>& verticesB, 
+	const b3AlignedObjectArray<b3Vector3>& uniqueEdgesB, 
+	const b3AlignedObjectArray<b3GpuFace>& facesB,
+	const b3AlignedObjectArray<int>& indicesB,
+
+	b3Vector3& sep)
+{
+	B3_PROFILE("findSeparatingAxis");
+
+	b3g_actualSATPairTests++;
+	float4 posA = posA1;
+	posA.w = 0.f;
+	float4 posB = posB1;
+	posB.w = 0.f;
+//#ifdef TEST_INTERNAL_OBJECTS
+	float4 c0local = (float4&)hullA.m_localCenter;
+	float4 c0 = transform(&c0local, &posA, &ornA);
+	float4 c1local = (float4&)hullB.m_localCenter;
+	float4 c1 = transform(&c1local,&posB,&ornB);
+	const float4 deltaC2 = c0 - c1;
+//#endif
+
+	b3Scalar dmin = FLT_MAX;
+	int curPlaneTests=0;
+
+	int numFacesA = hullA.m_numFaces;
+	// Test normals from hullA
+	for(int i=0;i<numFacesA;i++)
+	{
+		const float4& normal = (float4&)facesA[hullA.m_faceOffset+i].m_plane;
+		float4 faceANormalWS = b3QuatRotate(ornA,normal);
+
+		if (dot3F4(deltaC2,faceANormalWS)<0)
+			faceANormalWS*=-1.f;
+
+		curPlaneTests++;
+#ifdef TEST_INTERNAL_OBJECTS
+		gExpectedNbTests++;
+		if(gUseInternalObject && !TestInternalObjects(transA,transB, DeltaC2, faceANormalWS, hullA, hullB, dmin))
+			continue;
+		gActualNbTests++;
+#endif
+
+		
+		b3Scalar d;
+		if(!TestSepAxis( hullA, hullB, posA,ornA,posB,ornB,faceANormalWS, verticesA, verticesB,d))
+			return false;
+
+		if(d<dmin)
+		{
+			dmin = d;
+			sep = (b3Vector3&)faceANormalWS;
+		}
+	}
+
+	int numFacesB = hullB.m_numFaces;
+	// Test normals from hullB
+	for(int i=0;i<numFacesB;i++)
+	{
+		float4 normal = (float4&)facesB[hullB.m_faceOffset+i].m_plane;
+		float4 WorldNormal = b3QuatRotate(ornB, normal);
+
+		if (dot3F4(deltaC2,WorldNormal)<0)
+		{
+			WorldNormal*=-1.f;
+		}
+		curPlaneTests++;
+#ifdef TEST_INTERNAL_OBJECTS
+		gExpectedNbTests++;
+		if(gUseInternalObject && !TestInternalObjects(transA,transB,DeltaC2, WorldNormal, hullA, hullB, dmin))
+			continue;
+		gActualNbTests++;
+#endif
+
+		b3Scalar d;
+		if(!TestSepAxis(hullA, hullB,posA,ornA,posB,ornB,WorldNormal,verticesA,verticesB,d))
+			return false;
+
+		if(d<dmin)
+		{
+			dmin = d;
+			sep = (b3Vector3&)WorldNormal;
+		}
+	}
+
+	b3Vector3 edgeAstart,edgeAend,edgeBstart,edgeBend;
+
+	int curEdgeEdge = 0;
+	// Test edges
+	for(int e0=0;e0<hullA.m_numUniqueEdges;e0++)
+	{
+		const float4& edge0 = (float4&) uniqueEdgesA[hullA.m_uniqueEdgesOffset+e0];
+		float4 edge0World = b3QuatRotate(ornA,(float4&)edge0);
+
+		for(int e1=0;e1<hullB.m_numUniqueEdges;e1++)
+		{
+			const b3Vector3 edge1 = uniqueEdgesB[hullB.m_uniqueEdgesOffset+e1];
+			float4 edge1World = b3QuatRotate(ornB,(float4&)edge1);
+
+
+			float4 crossje = cross3(edge0World,edge1World);
+
+			curEdgeEdge++;
+			if(!IsAlmostZero((b3Vector3&)crossje))
+			{
+				crossje = normalize3(crossje);
+				if (dot3F4(deltaC2,crossje)<0)
+					crossje*=-1.f;
+
+
+#ifdef TEST_INTERNAL_OBJECTS
+				gExpectedNbTests++;
+				if(gUseInternalObject && !TestInternalObjects(transA,transB,DeltaC2, Cross, hullA, hullB, dmin))
+					continue;
+				gActualNbTests++;
+#endif
+
+				b3Scalar dist;
+				if(!TestSepAxis( hullA, hullB, posA,ornA,posB,ornB,crossje, verticesA,verticesB,dist))
+					return false;
+
+				if(dist<dmin)
+				{
+					dmin = dist;
+					sep = (b3Vector3&)crossje;
+				}
+			}
+		}
+
+	}
+
+	
+	if((dot3F4(-deltaC2,(float4&)sep))>0.0f)
+		sep = -sep;
+
+	return true;
+}
+
+
+bool findSeparatingAxisEdgeEdge(	__global const b3ConvexPolyhedronData* hullA, __global const b3ConvexPolyhedronData* hullB, 
+	const b3Float4& posA1,
+	const b3Quat& ornA,
+	const b3Float4& posB1,
+	const b3Quat& ornB,
+	const b3Float4& DeltaC2,
+	__global const b3AlignedObjectArray<float4>& vertices, 
+	__global const b3AlignedObjectArray<float4>& uniqueEdges, 
+	__global const b3AlignedObjectArray<b3GpuFace>& faces,
+	__global const b3AlignedObjectArray<int>&  indices,
+	float4* sep,
+	float* dmin)
+{
+//	int i = get_global_id(0);
+
+	float4 posA = posA1;
+	posA.w = 0.f;
+	float4 posB = posB1;
+	posB.w = 0.f;
+
+	int curPlaneTests=0;
+
+	int curEdgeEdge = 0;
+	// Test edges
+	for(int e0=0;e0<hullA->m_numUniqueEdges;e0++)
+	{
+		const float4 edge0 = uniqueEdges[hullA->m_uniqueEdgesOffset+e0];
+		float4 edge0World = b3QuatRotate(ornA,edge0);
+
+		for(int e1=0;e1<hullB->m_numUniqueEdges;e1++)
+		{
+			const float4 edge1 = uniqueEdges[hullB->m_uniqueEdgesOffset+e1];
+			float4 edge1World = b3QuatRotate(ornB,edge1);
+
+
+			float4 crossje = cross3(edge0World,edge1World);
+
+			curEdgeEdge++;
+			if(!IsAlmostZero(crossje))
+			{
+				crossje = normalize3(crossje);
+				if (dot3F4(DeltaC2,crossje)<0)
+					crossje*=-1.f;
+					
+				float dist;
+				bool result = true;
+				{
+					float Min0,Max0;
+					float Min1,Max1;
+					project(*hullA,posA,ornA,crossje,vertices, Min0, Max0);
+					project(*hullB,posB,ornB,crossje,vertices, Min1, Max1);
+				
+					if(Max0<Min1 || Max1<Min0)
+						result = false;
+				
+					float d0 = Max0 - Min1;
+					float d1 = Max1 - Min0;
+					dist = d0<d1 ? d0:d1;
+					result = true;
+
+				}
+				
+
+				if(dist<*dmin)
+				{
+					*dmin = dist;
+					*sep = crossje;
+				}
+			}
+		}
+
+	}
+
+	
+	if((dot3F4(-DeltaC2,*sep))>0.0f)
+	{
+		*sep = -(*sep);
+	}
+	return true;
+}
+
+
+__inline float4 lerp3(const float4& a,const float4& b, float  t)
+{
+	return make_float4(	a.x + (b.x - a.x) * t,
+						a.y + (b.y - a.y) * t,
+						a.z + (b.z - a.z) * t,
+						0.f);
+}
+
+
+// Clips a face to the back of a plane, return the number of vertices out, stored in ppVtxOut
+int clipFace(const float4* pVtxIn, int numVertsIn, float4& planeNormalWS,float planeEqWS, float4* ppVtxOut)
+{
+	
+	int ve;
+	float ds, de;
+	int numVertsOut = 0;
+	if (numVertsIn < 2)
+		return 0;
+
+	float4 firstVertex=pVtxIn[numVertsIn-1];
+	float4 endVertex = pVtxIn[0];
+	
+	ds = dot3F4(planeNormalWS,firstVertex)+planeEqWS;
+
+	for (ve = 0; ve < numVertsIn; ve++)
+	{
+		endVertex=pVtxIn[ve];
+
+		de = dot3F4(planeNormalWS,endVertex)+planeEqWS;
+
+		if (ds<0)
+		{
+			if (de<0)
+			{
+				// Start < 0, end < 0, so output endVertex
+				ppVtxOut[numVertsOut++] = endVertex;
+			}
+			else
+			{
+				// Start < 0, end >= 0, so output intersection
+				ppVtxOut[numVertsOut++] = lerp3(firstVertex, endVertex,(ds * 1.f/(ds - de)) );
+			}
+		}
+		else
+		{
+			if (de<0)
+			{
+				// Start >= 0, end < 0 so output intersection and end
+				ppVtxOut[numVertsOut++] = lerp3(firstVertex, endVertex,(ds * 1.f/(ds - de)) );
+				ppVtxOut[numVertsOut++] = endVertex;
+			}
+		}
+		firstVertex = endVertex;
+		ds = de;
+	}
+	return numVertsOut;
+}
+
+
+int clipFaceAgainstHull(const float4& separatingNormal, const b3ConvexPolyhedronData* hullA,  
+	const float4& posA, const b3Quaternion& ornA, float4* worldVertsB1, int numWorldVertsB1,
+	float4* worldVertsB2, int capacityWorldVertsB2,
+	const float minDist, float maxDist,
+	const b3AlignedObjectArray<float4>& verticesA,	const b3AlignedObjectArray<b3GpuFace>& facesA,	const b3AlignedObjectArray<int>& indicesA,
+	//const float4* verticesB,	const b3GpuFace* facesB,	const int* indicesB,
+	float4* contactsOut,
+	int contactCapacity)
+{
+	int numContactsOut = 0;
+
+	float4* pVtxIn = worldVertsB1;
+	float4* pVtxOut = worldVertsB2;
+	
+	int numVertsIn = numWorldVertsB1;
+	int numVertsOut = 0;
+
+	int closestFaceA=-1;
+	{
+		float dmin = FLT_MAX;
+		for(int face=0;face<hullA->m_numFaces;face++)
+		{
+			const float4 Normal = make_float4(
+				facesA[hullA->m_faceOffset+face].m_plane.x, 
+				facesA[hullA->m_faceOffset+face].m_plane.y, 
+				facesA[hullA->m_faceOffset+face].m_plane.z,0.f);
+			const float4 faceANormalWS = b3QuatRotate(ornA,Normal);
+		
+			float d = dot3F4(faceANormalWS,separatingNormal);
+			if (d < dmin)
+			{
+				dmin = d;
+				closestFaceA = face;
+			}
+		}
+	}
+	if (closestFaceA<0)
+		return numContactsOut;
+
+	b3GpuFace polyA = facesA[hullA->m_faceOffset+closestFaceA];
+
+	// clip polygon to back of planes of all faces of hull A that are adjacent to witness face
+	int numContacts = numWorldVertsB1;
+	int numVerticesA = polyA.m_numIndices;
+	for(int e0=0;e0<numVerticesA;e0++)
+	{
+		const float4 a = verticesA[hullA->m_vertexOffset+indicesA[polyA.m_indexOffset+e0]];
+		const float4 b = verticesA[hullA->m_vertexOffset+indicesA[polyA.m_indexOffset+((e0+1)%numVerticesA)]];
+		const float4 edge0 = a - b;
+		const float4 WorldEdge0 = b3QuatRotate(ornA,edge0);
+		float4 planeNormalA = make_float4(polyA.m_plane.x,polyA.m_plane.y,polyA.m_plane.z,0.f);
+		float4 worldPlaneAnormal1 = b3QuatRotate(ornA,planeNormalA);
+
+		float4 planeNormalWS1 = -cross3(WorldEdge0,worldPlaneAnormal1);
+		float4 worldA1 = transform(&a,&posA,&ornA);
+		float planeEqWS1 = -dot3F4(worldA1,planeNormalWS1);
+		
+		float4 planeNormalWS = planeNormalWS1;
+		float planeEqWS=planeEqWS1;
+		
+		//clip face
+		//clipFace(*pVtxIn, *pVtxOut,planeNormalWS,planeEqWS);
+		numVertsOut = clipFace(pVtxIn, numVertsIn, planeNormalWS,planeEqWS, pVtxOut);
+
+		//btSwap(pVtxIn,pVtxOut);
+		float4* tmp = pVtxOut;
+		pVtxOut = pVtxIn;
+		pVtxIn = tmp;
+		numVertsIn = numVertsOut;
+		numVertsOut = 0;
+	}
+
+	
+	// only keep points that are behind the witness face
+	{
+		float4 localPlaneNormal  = make_float4(polyA.m_plane.x,polyA.m_plane.y,polyA.m_plane.z,0.f);
+		float localPlaneEq = polyA.m_plane.w;
+		float4 planeNormalWS = b3QuatRotate(ornA,localPlaneNormal);
+		float planeEqWS=localPlaneEq-dot3F4(planeNormalWS,posA);
+		for (int i=0;i<numVertsIn;i++)
+		{
+			float depth = dot3F4(planeNormalWS,pVtxIn[i])+planeEqWS;
+			if (depth <=minDist)
+			{
+				depth = minDist;
+			}
+			if (numContactsOut<contactCapacity)
+			{
+				if (depth <=maxDist)
+				{
+					float4 pointInWorld = pVtxIn[i];
+					//resultOut.addContactPoint(separatingNormal,point,depth);
+					contactsOut[numContactsOut++] = make_float4(pointInWorld.x,pointInWorld.y,pointInWorld.z,depth);
+					//printf("depth=%f\n",depth);
+				}
+			} else
+			{
+				b3Error("exceeding contact capacity (%d,%df)\n", numContactsOut,contactCapacity);
+			}
+		}
+	}
+
+	return numContactsOut;
+}
+
+
+
+static int	clipHullAgainstHull(const float4& separatingNormal, 
+	const b3ConvexPolyhedronData& hullA, const b3ConvexPolyhedronData& hullB, 
+	const float4& posA, const b3Quaternion& ornA,const float4& posB, const b3Quaternion& ornB, 
+	float4* worldVertsB1, float4* worldVertsB2, int capacityWorldVerts,
+	const float minDist, float maxDist,
+	const b3AlignedObjectArray<float4>& verticesA,	const b3AlignedObjectArray<b3GpuFace>& facesA,	const b3AlignedObjectArray<int>& indicesA,
+	const b3AlignedObjectArray<float4>& verticesB,	const b3AlignedObjectArray<b3GpuFace>& facesB,	const b3AlignedObjectArray<int>& indicesB,
+
+	float4*	contactsOut,
+	int contactCapacity)
+{
+	int numContactsOut = 0;
+	int numWorldVertsB1= 0;
+	
+	B3_PROFILE("clipHullAgainstHull");
+
+	float curMaxDist=maxDist;
+	int closestFaceB=-1;
+	float dmax = -FLT_MAX;
+
+	{
+		//B3_PROFILE("closestFaceB");
+		if (hullB.m_numFaces!=1)
+		{
+			//printf("wtf\n");
+		}
+		static bool once = true;
+		//printf("separatingNormal=%f,%f,%f\n",separatingNormal.x,separatingNormal.y,separatingNormal.z);
+		
+		for(int face=0;face<hullB.m_numFaces;face++)
+		{
+#ifdef BT_DEBUG_SAT_FACE
+			if (once)
+				printf("face %d\n",face);
+			const b3GpuFace* faceB = &facesB[hullB.m_faceOffset+face];
+			if (once)
+			{
+				for (int i=0;i<faceB->m_numIndices;i++)
+				{
+					float4 vert = verticesB[hullB.m_vertexOffset+indicesB[faceB->m_indexOffset+i]];
+					printf("vert[%d] = %f,%f,%f\n",i,vert.x,vert.y,vert.z);
+				}
+			}
+#endif //BT_DEBUG_SAT_FACE
+			//if (facesB[hullB.m_faceOffset+face].m_numIndices>2)
+			{
+				const float4 Normal = make_float4(facesB[hullB.m_faceOffset+face].m_plane.x, 
+					facesB[hullB.m_faceOffset+face].m_plane.y, facesB[hullB.m_faceOffset+face].m_plane.z,0.f);
+				const float4 WorldNormal = b3QuatRotate(ornB, Normal);
+#ifdef BT_DEBUG_SAT_FACE
+				if (once)
+					printf("faceNormal = %f,%f,%f\n",Normal.x,Normal.y,Normal.z);
+#endif
+				float d = dot3F4(WorldNormal,separatingNormal);
+				if (d > dmax)
+				{
+					dmax = d;
+					closestFaceB = face;
+				}
+			}
+		}
+		once = false;
+	}
+
+	
+	b3Assert(closestFaceB>=0);
+	{
+		//B3_PROFILE("worldVertsB1");
+		const b3GpuFace& polyB = facesB[hullB.m_faceOffset+closestFaceB];
+		const int numVertices = polyB.m_numIndices;
+		for(int e0=0;e0<numVertices;e0++)
+		{
+			const float4& b = verticesB[hullB.m_vertexOffset+indicesB[polyB.m_indexOffset+e0]];
+			worldVertsB1[numWorldVertsB1++] = transform(&b,&posB,&ornB);
+		}
+	}
+
+	if (closestFaceB>=0)
+	{
+		//B3_PROFILE("clipFaceAgainstHull");
+		numContactsOut = clipFaceAgainstHull((float4&)separatingNormal, &hullA, 
+				posA,ornA,
+				worldVertsB1,numWorldVertsB1,worldVertsB2,capacityWorldVerts, minDist, maxDist,
+				verticesA,				facesA,				indicesA,
+				contactsOut,contactCapacity);
+	}
+
+	return numContactsOut;
+}
+
+
+
+
+
+
+#define PARALLEL_SUM(v, n) for(int j=1; j<n; j++) v[0] += v[j];
+#define PARALLEL_DO(execution, n) for(int ie=0; ie<n; ie++){execution;}
+#define REDUCE_MAX(v, n) {int i=0;\
+for(int offset=0; offset<n; offset++) v[i] = (v[i].y > v[i+offset].y)? v[i]: v[i+offset]; }
+#define REDUCE_MIN(v, n) {int i=0;\
+for(int offset=0; offset<n; offset++) v[i] = (v[i].y < v[i+offset].y)? v[i]: v[i+offset]; }
+
+int extractManifold(const float4* p, int nPoints, const float4& nearNormal, b3Int4* contactIdx)
+{
+	if( nPoints == 0 )
+        return 0;
+    
+    if (nPoints <=4)
+        return nPoints;
+    
+    
+    if (nPoints >64)
+        nPoints = 64;
+    
+	float4 center = make_float4(0,0,0,0);
+	{
+		
+		for (int i=0;i<nPoints;i++)
+			center += p[i];
+		center /= (float)nPoints;
+	}
+    
+	
+    
+	//	sample 4 directions
+    
+    float4 aVector = p[0] - center;
+    float4 u = cross3( nearNormal, aVector );
+    float4 v = cross3( nearNormal, u );
+    u = normalize3( u );
+    v = normalize3( v );
+    
+    
+    //keep point with deepest penetration
+    float minW= FLT_MAX;
+    
+    int minIndex=-1;
+    
+    float4 maxDots;
+    maxDots.x = FLT_MIN;
+    maxDots.y = FLT_MIN;
+    maxDots.z = FLT_MIN;
+    maxDots.w = FLT_MIN;
+    
+    //	idx, distance
+    for(int ie = 0; ie<nPoints; ie++ )
+    {
+        if (p[ie].w<minW)
+        {
+            minW = p[ie].w;
+            minIndex=ie;
+        }
+        float f;
+        float4 r = p[ie]-center;
+        f = dot3F4( u, r );
+        if (f<maxDots.x)
+        {
+            maxDots.x = f;
+            contactIdx[0].x = ie;
+        }
+        
+        f = dot3F4( -u, r );
+        if (f<maxDots.y)
+        {
+            maxDots.y = f;
+            contactIdx[0].y = ie;
+        }
+        
+        
+        f = dot3F4( v, r );
+        if (f<maxDots.z)
+        {
+            maxDots.z = f;
+            contactIdx[0].z = ie;
+        }
+        
+        f = dot3F4( -v, r );
+        if (f<maxDots.w)
+        {
+            maxDots.w = f;
+            contactIdx[0].w = ie;
+        }
+        
+    }
+    
+    if (contactIdx[0].x != minIndex && contactIdx[0].y != minIndex && contactIdx[0].z != minIndex && contactIdx[0].w != minIndex)
+    {
+        //replace the first contact with minimum (todo: replace contact with least penetration)
+        contactIdx[0].x = minIndex;
+    }
+    
+    return 4;
+    
+}
+
+
+
+
+int clipHullHullSingle(
+			int bodyIndexA, int bodyIndexB,
+										 const float4& posA,
+										 const b3Quaternion& ornA,
+										 const float4& posB,
+										 const b3Quaternion& ornB,
+
+			int collidableIndexA, int collidableIndexB,
+
+			const b3AlignedObjectArray<b3RigidBodyCL>* bodyBuf, 
+			b3AlignedObjectArray<b3Contact4>* globalContactOut, 
+			int& nContacts,
+			
+			const b3AlignedObjectArray<b3ConvexPolyhedronCL>& hostConvexDataA,
+			const b3AlignedObjectArray<b3ConvexPolyhedronCL>& hostConvexDataB,
+	
+			const b3AlignedObjectArray<b3Vector3>& verticesA, 
+			const b3AlignedObjectArray<b3Vector3>& uniqueEdgesA, 
+			const b3AlignedObjectArray<b3GpuFace>& facesA,
+			const b3AlignedObjectArray<int>& indicesA,
+	
+			const b3AlignedObjectArray<b3Vector3>& verticesB,
+			const b3AlignedObjectArray<b3Vector3>& uniqueEdgesB,
+			const b3AlignedObjectArray<b3GpuFace>& facesB,
+			const b3AlignedObjectArray<int>& indicesB,
+
+			const b3AlignedObjectArray<b3Collidable>& hostCollidablesA,
+			const b3AlignedObjectArray<b3Collidable>& hostCollidablesB,
+			const b3Vector3& sepNormalWorldSpace,
+			int maxContactCapacity			)
+{
+	int contactIndex = -1;
+	b3ConvexPolyhedronCL hullA, hullB;
+    
+    b3Collidable colA = hostCollidablesA[collidableIndexA];
+    hullA = hostConvexDataA[colA.m_shapeIndex];
+    //printf("numvertsA = %d\n",hullA.m_numVertices);
+    
+    
+    b3Collidable colB = hostCollidablesB[collidableIndexB];
+    hullB = hostConvexDataB[colB.m_shapeIndex];
+    //printf("numvertsB = %d\n",hullB.m_numVertices);
+    
+	
+	float4 contactsOut[MAX_VERTS];
+	int localContactCapacity = MAX_VERTS;
+
+#ifdef _WIN32
+	b3Assert(_finite(bodyBuf->at(bodyIndexA).m_pos.x));
+	b3Assert(_finite(bodyBuf->at(bodyIndexB).m_pos.x));
+#endif
+	
+	
+	{
+		
+		float4 worldVertsB1[MAX_VERTS];
+		float4 worldVertsB2[MAX_VERTS];
+		int capacityWorldVerts = MAX_VERTS;
+
+		float4 hostNormal = make_float4(sepNormalWorldSpace.getX(),sepNormalWorldSpace.getY(),sepNormalWorldSpace.getZ(),0.f);
+		int shapeA = hostCollidablesA[collidableIndexA].m_shapeIndex;
+		int shapeB = hostCollidablesB[collidableIndexB].m_shapeIndex;
+
+		b3Scalar minDist = -1;
+		b3Scalar maxDist = 0.;
+
+		        
+
+		b3Transform trA,trB;
+		{
+		//B3_PROFILE("transform computation");
+		//trA.setIdentity();
+		trA.setOrigin(b3Vector3(posA.x,posA.y,posA.z));
+		trA.setRotation(b3Quaternion(ornA.x,ornA.y,ornA.z,ornA.w));
+				
+		//trB.setIdentity();
+		trB.setOrigin(b3Vector3(posB.x,posB.y,posB.z));
+		trB.setRotation(b3Quaternion(ornB.x,ornB.y,ornB.z,ornB.w));
+		}
+
+		b3Quaternion trAorn = trA.getRotation();
+        b3Quaternion trBorn = trB.getRotation();
+        
+		int numContactsOut = clipHullAgainstHull(hostNormal, 
+						hostConvexDataA.at(shapeA), 
+						hostConvexDataB.at(shapeB),
+								(float4&)trA.getOrigin(), (b3Quaternion&)trAorn,
+								(float4&)trB.getOrigin(), (b3Quaternion&)trBorn,
+								worldVertsB1,worldVertsB2,capacityWorldVerts,
+								minDist, maxDist,
+								verticesA,	facesA,indicesA,
+								verticesB,	facesB,indicesB,
+								
+								contactsOut,localContactCapacity);
+
+		if (numContactsOut>0)
+		{
+			B3_PROFILE("overlap");
+
+			float4 normalOnSurfaceB = (float4&)hostNormal;
+			float4 centerOut;
+			
+			b3Int4 contactIdx;
+			contactIdx.x = 0;
+			contactIdx.y = 1;
+			contactIdx.z = 2;
+			contactIdx.w = 3;
+			
+			int numPoints = 0;
+					
+			{
+				B3_PROFILE("extractManifold");
+				numPoints = extractManifold(contactsOut, numContactsOut, normalOnSurfaceB,  &contactIdx);
+			}
+					
+			b3Assert(numPoints);
+					
+			if (nContacts<maxContactCapacity)
+			{
+				contactIndex = nContacts;
+				globalContactOut->expand();
+				b3Contact4& contact = globalContactOut->at(nContacts);
+				contact.m_batchIdx = 0;//i;
+				contact.m_bodyAPtrAndSignBit = (bodyBuf->at(bodyIndexA).m_invMass==0)? -bodyIndexA:bodyIndexA;
+				contact.m_bodyBPtrAndSignBit = (bodyBuf->at(bodyIndexB).m_invMass==0)? -bodyIndexB:bodyIndexB;
+
+				contact.m_frictionCoeffCmp = 45874;
+				contact.m_restituitionCoeffCmp = 0;
+					
+				float distance = 0.f;
+				for (int p=0;p<numPoints;p++)
+				{
+					contact.m_worldPosB[p] = contactsOut[contactIdx.s[p]];//check if it is actually on B
+					contact.m_worldNormalOnB = normalOnSurfaceB; 
+				}
+				//printf("bodyIndexA %d,bodyIndexB %d,normal=%f,%f,%f numPoints %d\n",bodyIndexA,bodyIndexB,normalOnSurfaceB.x,normalOnSurfaceB.y,normalOnSurfaceB.z,numPoints);
+				contact.m_worldNormalOnB.w = (b3Scalar)numPoints;
+				nContacts++;
+			} else
+			{
+				b3Error("Error: exceeding contact capacity (%d/%d)\n", nContacts,maxContactCapacity);
+			}
+		}
+	}
+	return contactIndex;
+}
+
 	
 
 
@@ -557,6 +1368,617 @@ void computeContactPlaneConvex(int pairIndex,
 
 
 
+// work-in-progress
+__kernel void   findCompoundPairsKernel( 
+	int pairIndex,
+	int bodyIndexA,
+	int bodyIndexB,
+	int collidableIndexA,
+	int collidableIndexB,
+	__global const b3RigidBodyData* rigidBodies, 
+	__global const b3Collidable* collidables,
+	__global const b3ConvexPolyhedronData* convexShapes, 
+	__global const b3AlignedObjectArray<b3Float4>& vertices,
+	__global const b3AlignedObjectArray<b3Aabb>& aabbsWorldSpace,
+	__global const b3AlignedObjectArray<b3Aabb>& aabbsLocalSpace,
+	__global const b3GpuChildShape* gpuChildShapes,
+	__global b3Int4* gpuCompoundPairsOut,
+	__global  int* numCompoundPairsOut,
+	int maxNumCompoundPairsCapacity
+	)
+{
+
+
+	int i = pairIndex;
+	{
+		
+
+		int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;
+		int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;
+
+
+		//once the broadphase avoids static-static pairs, we can remove this test
+		if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))
+		{
+			return;
+		}
+
+		if ((collidables[collidableIndexA].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) ||(collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS))
+		{
+
+			if (collidables[collidableIndexA].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) 
+			{
+
+				int numChildrenA = collidables[collidableIndexA].m_numChildShapes;
+				for (int c=0;c<numChildrenA;c++)
+				{
+					int childShapeIndexA = collidables[collidableIndexA].m_shapeIndex+c;
+					int childColIndexA = gpuChildShapes[childShapeIndexA].m_shapeIndex;
+
+					float4 posA = rigidBodies[bodyIndexA].m_pos;
+					b3Quat ornA = rigidBodies[bodyIndexA].m_quat;
+					float4 childPosA = gpuChildShapes[childShapeIndexA].m_childPosition;
+					b3Quat childOrnA = gpuChildShapes[childShapeIndexA].m_childOrientation;
+					float4 newPosA = b3QuatRotate(ornA,childPosA)+posA;
+					b3Quat newOrnA = b3QuatMul(ornA,childOrnA);
+					
+
+					int shapeIndexA = collidables[childColIndexA].m_shapeIndex;
+					b3Aabb aabbA = aabbsLocalSpace[shapeIndexA];
+
+					
+					b3Transform transA;
+					transA.setIdentity();
+					transA.setOrigin(newPosA);
+					transA.setRotation(newOrnA);
+					b3Scalar margin=0.0f;
+
+					b3Vector3 aabbAMinOut,aabbAMaxOut;
+
+					b3TransformAabb((const b3Float4&)aabbA.m_min,(const b3Float4&)aabbA.m_max, margin,transA,aabbAMinOut,aabbAMaxOut);
+
+					if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)
+					{
+						int numChildrenB = collidables[collidableIndexB].m_numChildShapes;
+						for (int b=0;b<numChildrenB;b++)
+						{
+							int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+b;
+							int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;
+							b3Quat ornB = rigidBodies[bodyIndexB].m_quat;
+							float4 posB = rigidBodies[bodyIndexB].m_pos;
+							float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;
+							b3Quat childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;
+							float4 newPosB = transform(&childPosB,&posB,&ornB);
+							b3Quat newOrnB = b3QuatMul(ornB,childOrnB);
+
+							
+
+
+							int shapeIndexB = collidables[childColIndexB].m_shapeIndex;
+
+							b3Aabb aabbB = aabbsLocalSpace[shapeIndexB];
+
+							b3Transform transB;
+							transB.setIdentity();
+							transB.setOrigin(newPosB);
+							transB.setRotation(newOrnB);
+
+							b3Vector3 aabbBMinOut,aabbBMaxOut;
+							b3TransformAabb((const b3Float4&)aabbB.m_min,(const b3Float4&)aabbB.m_max, margin,transB,aabbBMinOut,aabbBMaxOut);
+
+							bool aabbOverlap = b3TestAabbAgainstAabb2(aabbAMinOut,aabbAMaxOut,aabbBMinOut,aabbBMaxOut);
+							if (aabbOverlap)
+							{
+								/*
+								int numFacesA = convexShapes[shapeIndexA].m_numFaces;
+								float dmin = FLT_MAX;
+								float4 posA = newPosA;
+								posA.w = 0.f;
+								float4 posB = newPosB;
+								posB.w = 0.f;
+								float4 c0local = convexShapes[shapeIndexA].m_localCenter;
+								b3Quat ornA = newOrnA;
+								float4 c0 = transform(&c0local, &posA, &ornA);
+								float4 c1local = convexShapes[shapeIndexB].m_localCenter;
+								b3Quat ornB =newOrnB;
+								float4 c1 = transform(&c1local,&posB,&ornB);
+								const float4 DeltaC2 = c0 - c1;
+								*/
+								{//
+									int compoundPairIdx = b3AtomicInc(numCompoundPairsOut);
+									if (compoundPairIdx<maxNumCompoundPairsCapacity)
+									{
+										gpuCompoundPairsOut[compoundPairIdx]  = b3MakeInt4(bodyIndexA,bodyIndexB,childShapeIndexA,childShapeIndexB);
+									}
+								}//
+							}//fi(1)
+						} //for (int b=0
+					}//if (collidables[collidableIndexB].
+					else//if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)
+					{
+						if (1)
+						{
+							int numFacesA = convexShapes[shapeIndexA].m_numFaces;
+							float dmin = FLT_MAX;
+							float4 posA = newPosA;
+							posA.w = 0.f;
+							float4 posB = rigidBodies[bodyIndexB].m_pos;
+							posB.w = 0.f;
+							float4 c0local = convexShapes[shapeIndexA].m_localCenter;
+							b3Quat ornA = newOrnA;
+							float4 c0 = transform(&c0local, &posA, &ornA);
+							float4 c1local = convexShapes[shapeIndexB].m_localCenter;
+							b3Quat ornB = rigidBodies[bodyIndexB].m_quat;
+							float4 c1 = transform(&c1local,&posB,&ornB);
+							const float4 DeltaC2 = c0 - c1;
+
+							{
+								int compoundPairIdx = b3AtomicInc(numCompoundPairsOut);
+								if (compoundPairIdx<maxNumCompoundPairsCapacity)
+								{
+									gpuCompoundPairsOut[compoundPairIdx] = b3MakeInt4(bodyIndexA,bodyIndexB,childShapeIndexA,-1);
+								}//if (compoundPairIdx<maxNumCompoundPairsCapacity)
+							}//
+						}//fi (1)
+					}//if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)
+				}//for (int b=0;b<numChildrenB;b++)	
+				return;
+			}//if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)
+			if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONCAVE_TRIMESH) 
+				&& (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS))
+			{
+				int numChildrenB = collidables[collidableIndexB].m_numChildShapes;
+				for (int b=0;b<numChildrenB;b++)
+				{
+					int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+b;
+					int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;
+					b3Quat ornB = rigidBodies[bodyIndexB].m_quat;
+					float4 posB = rigidBodies[bodyIndexB].m_pos;
+					float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;
+					b3Quat childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;
+					float4 newPosB = b3QuatRotate(ornB,childPosB)+posB;
+					b3Quat newOrnB = b3QuatMul(ornB,childOrnB);
+
+					int shapeIndexB = collidables[childColIndexB].m_shapeIndex;
+
+
+					//////////////////////////////////////
+
+					if (1)
+					{
+						int numFacesA = convexShapes[shapeIndexA].m_numFaces;
+						float dmin = FLT_MAX;
+						float4 posA = rigidBodies[bodyIndexA].m_pos;
+						posA.w = 0.f;
+						float4 posB = newPosB;
+						posB.w = 0.f;
+						float4 c0local = convexShapes[shapeIndexA].m_localCenter;
+						b3Quat ornA = rigidBodies[bodyIndexA].m_quat;
+						float4 c0 = transform(&c0local, &posA, &ornA);
+						float4 c1local = convexShapes[shapeIndexB].m_localCenter;
+						b3Quat ornB =newOrnB;
+						float4 c1 = transform(&c1local,&posB,&ornB);
+						const float4 DeltaC2 = c0 - c1;
+						{//
+							int compoundPairIdx = b3AtomicInc(numCompoundPairsOut);
+							if (compoundPairIdx<maxNumCompoundPairsCapacity)
+							{
+								gpuCompoundPairsOut[compoundPairIdx] = b3MakeInt4(bodyIndexA,bodyIndexB,-1,childShapeIndexB);
+							}//fi (compoundPairIdx<maxNumCompoundPairsCapacity)
+						}//
+					}//fi (1)	
+				}//for (int b=0;b<numChildrenB;b++)
+				return;
+			}//if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)
+			return;
+		}//fi ((collidables[collidableIndexA].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) ||(collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS))
+	}//i<numPairs
+}
+
+
+
+__kernel void   processCompoundPairsKernel( __global const b3Int4* gpuCompoundPairs,
+										__global const b3RigidBodyData* rigidBodies, 
+										__global const b3Collidable* collidables,
+										__global const b3ConvexPolyhedronData* convexShapes, 
+										__global const b3AlignedObjectArray<b3Float4>& vertices,
+										__global const b3AlignedObjectArray<b3Float4>& uniqueEdges,
+										__global const b3AlignedObjectArray<b3GpuFace>& faces,
+										__global const b3AlignedObjectArray<int>& indices,
+										__global b3Aabb* aabbs,
+										__global const b3GpuChildShape* gpuChildShapes,
+										__global b3AlignedObjectArray<b3Float4>& gpuCompoundSepNormalsOut,
+										__global b3AlignedObjectArray<int>& gpuHasCompoundSepNormalsOut,
+										int numCompoundPairs,
+										int i
+										)
+{
+
+//	int i = get_global_id(0);
+	if (i<numCompoundPairs)
+	{
+		int bodyIndexA = gpuCompoundPairs[i].x;
+		int bodyIndexB = gpuCompoundPairs[i].y;
+
+		int childShapeIndexA = gpuCompoundPairs[i].z;
+		int childShapeIndexB = gpuCompoundPairs[i].w;
+		
+		int collidableIndexA = -1;
+		int collidableIndexB = -1;
+		
+		b3Quat ornA = rigidBodies[bodyIndexA].m_quat;
+		float4 posA = rigidBodies[bodyIndexA].m_pos;
+		
+		b3Quat ornB = rigidBodies[bodyIndexB].m_quat;
+		float4 posB = rigidBodies[bodyIndexB].m_pos;
+							
+		if (childShapeIndexA >= 0)
+		{
+			collidableIndexA = gpuChildShapes[childShapeIndexA].m_shapeIndex;
+			float4 childPosA = gpuChildShapes[childShapeIndexA].m_childPosition;
+			b3Quat	childOrnA = gpuChildShapes[childShapeIndexA].m_childOrientation;
+			float4 newPosA = b3QuatRotate(ornA,childPosA)+posA;
+			b3Quat newOrnA = b3QuatMul(ornA,childOrnA);
+			posA = newPosA;
+			ornA = newOrnA;
+		} else
+		{
+			collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;
+		}
+		
+		if (childShapeIndexB>=0)
+		{
+			collidableIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;
+			float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;
+			b3Quat childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;
+			float4 newPosB = b3QuatRotate(ornB,childPosB)+posB;
+			b3Quat newOrnB = b3QuatMul(ornB,childOrnB);
+			posB = newPosB;
+			ornB = newOrnB;
+		} else
+		{
+			collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;	
+		}
+	
+		gpuHasCompoundSepNormalsOut[i] = 0;
+	
+		int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;
+		int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;
+	
+		int shapeTypeA = collidables[collidableIndexA].m_shapeType;
+		int shapeTypeB = collidables[collidableIndexB].m_shapeType;
+	
+
+		if ((shapeTypeA != SHAPE_CONVEX_HULL) || (shapeTypeB != SHAPE_CONVEX_HULL))
+		{
+			return;
+		}
+
+		int hasSeparatingAxis = 5;
+							
+		int numFacesA = convexShapes[shapeIndexA].m_numFaces;
+		float dmin = FLT_MAX;
+		posA.w = 0.f;
+		posB.w = 0.f;
+		float4 c0local = convexShapes[shapeIndexA].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;
+		float4 sepNormal = make_float4(1,0,0,0);
+//		bool sepA = findSeparatingAxis(	convexShapes[shapeIndexA], convexShapes[shapeIndexB],posA,ornA,posB,ornB,DeltaC2,vertices,uniqueEdges,faces,indices,&sepNormal,&dmin);
+		bool sepA = findSeparatingAxis(	convexShapes[shapeIndexA], convexShapes[shapeIndexB],posA,ornA,posB,ornB,vertices,uniqueEdges,faces,indices,vertices,uniqueEdges,faces,indices,sepNormal);//,&dmin);
+	
+		hasSeparatingAxis = 4;
+		if (!sepA)
+		{
+			hasSeparatingAxis = 0;
+		} else
+		{
+			bool sepB = findSeparatingAxis(	convexShapes[shapeIndexB],convexShapes[shapeIndexA],posB,ornB,posA,ornA,vertices,uniqueEdges,faces,indices,vertices,uniqueEdges,faces,indices,sepNormal);//,&dmin);
+
+			if (!sepB)
+			{
+				hasSeparatingAxis = 0;
+			} else//(!sepB)
+			{
+				bool sepEE = findSeparatingAxisEdgeEdge(	&convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA,posB,ornB,DeltaC2,vertices,uniqueEdges,faces,indices,&sepNormal,&dmin);
+				if (sepEE)
+				{
+						gpuCompoundSepNormalsOut[i] = sepNormal;//fastNormalize4(sepNormal);
+						gpuHasCompoundSepNormalsOut[i] = 1;
+				}//sepEE
+			}//(!sepB)
+		}//(!sepA)
+		
+		
+	}
+		
+}
+
+
+__kernel void   clipCompoundsHullHullKernel( __global const b3Int4* gpuCompoundPairs, 
+																					__global const b3RigidBodyData* rigidBodies, 
+																					__global const b3Collidable* collidables,
+																					__global const b3ConvexPolyhedronData* convexShapes, 
+																					__global const b3AlignedObjectArray<b3Float4>& vertices,
+																					__global const b3AlignedObjectArray<b3Float4>& uniqueEdges,
+																					__global const b3AlignedObjectArray<b3GpuFace>& faces,
+																					__global const b3AlignedObjectArray<int>& indices,
+																					__global const b3GpuChildShape* gpuChildShapes,
+																					__global const b3AlignedObjectArray<b3Float4>& gpuCompoundSepNormalsOut,
+																					__global const b3AlignedObjectArray<int>& gpuHasCompoundSepNormalsOut,
+																					__global struct b3Contact4Data* globalContactsOut,
+																					int* nGlobalContactsOut,
+																					int numCompoundPairs, int maxContactCapacity, int i)
+{
+
+//	int i = get_global_id(0);
+	int pairIndex = i;
+	
+	float4 worldVertsB1[64];
+	float4 worldVertsB2[64];
+	int capacityWorldVerts = 64;	
+
+	float4 localContactsOut[64];
+	int localContactCapacity=64;
+	
+	float minDist = -1e30f;
+	float maxDist = 0.0f;
+
+	if (i<numCompoundPairs)
+	{
+
+		if (gpuHasCompoundSepNormalsOut[i])
+		{
+
+			int bodyIndexA = gpuCompoundPairs[i].x;
+			int bodyIndexB = gpuCompoundPairs[i].y;
+			
+			int childShapeIndexA = gpuCompoundPairs[i].z;
+			int childShapeIndexB = gpuCompoundPairs[i].w;
+			
+			int collidableIndexA = -1;
+			int collidableIndexB = -1;
+			
+			b3Quat ornA = rigidBodies[bodyIndexA].m_quat;
+			float4 posA = rigidBodies[bodyIndexA].m_pos;
+			
+			b3Quat ornB = rigidBodies[bodyIndexB].m_quat;
+			float4 posB = rigidBodies[bodyIndexB].m_pos;
+								
+			if (childShapeIndexA >= 0)
+			{
+				collidableIndexA = gpuChildShapes[childShapeIndexA].m_shapeIndex;
+				float4 childPosA = gpuChildShapes[childShapeIndexA].m_childPosition;
+				b3Quat childOrnA = gpuChildShapes[childShapeIndexA].m_childOrientation;
+				float4 newPosA = b3QuatRotate(ornA,childPosA)+posA;
+				b3Quat newOrnA = b3QuatMul(ornA,childOrnA);
+				posA = newPosA;
+				ornA = newOrnA;
+			} else
+			{
+				collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;
+			}
+			
+			if (childShapeIndexB>=0)
+			{
+				collidableIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;
+				float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;
+				b3Quat  childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;
+				float4 newPosB = b3QuatRotate(ornB,childPosB)+posB;
+				b3Quat  newOrnB = b3QuatMul(ornB,childOrnB);
+				posB = newPosB;
+				ornB = newOrnB;
+			} else
+			{
+				collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;	
+			}
+			
+			int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;
+			int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;
+		
+			int numLocalContactsOut = clipHullAgainstHull(gpuCompoundSepNormalsOut[i],
+														convexShapes[shapeIndexA], convexShapes[shapeIndexB],
+														posA,ornA,
+													  posB,ornB,
+													  worldVertsB1,worldVertsB2,capacityWorldVerts,
+														minDist, maxDist,
+														vertices,faces,indices,
+														vertices,faces,indices,
+														localContactsOut,localContactCapacity);
+												
+		if (numLocalContactsOut>0)
+		{
+				float4 normal = -gpuCompoundSepNormalsOut[i];
+				int nPoints = numLocalContactsOut;
+				float4* pointsIn = localContactsOut;
+				b3Int4 contactIdx;// = {-1,-1,-1,-1};
+
+				contactIdx.s[0] = 0;
+				contactIdx.s[1] = 1;
+				contactIdx.s[2] = 2;
+				contactIdx.s[3] = 3;
+
+				int nReducedContacts = extractManifoldSequentialGlobal(pointsIn, nPoints, normal, &contactIdx);
+		
+				int dstIdx;
+				dstIdx = b3AtomicInc( nGlobalContactsOut);
+				if ((dstIdx+nReducedContacts) < maxContactCapacity)
+				{
+					__global struct b3Contact4Data* c = globalContactsOut+ dstIdx;
+					c->m_worldNormalOnB = -normal;
+					c->m_restituitionCoeffCmp = (0.f*0xffff);c->m_frictionCoeffCmp = (0.7f*0xffff);
+					c->m_batchIdx = pairIndex;
+					int bodyA = gpuCompoundPairs[pairIndex].x;
+					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_worldPosB[i] = pointsIn[contactIdx.s[i]];
+					}
+					b3Contact4Data_setNumPoints(c,nReducedContacts);
+				}
+				
+			}//		if (numContactsOut>0)
+		}//		if (gpuHasCompoundSepNormalsOut[i])
+	}//	if (i<numCompoundPairs)
+
+}
+
+
+void computeContactCompoundCompound(int pairIndex,
+																int bodyIndexA, int bodyIndexB, 
+																int collidableIndexA, int collidableIndexB, 
+																const b3RigidBodyCL* rigidBodies, 
+																const b3Collidable* collidables,
+																const b3ConvexPolyhedronData* convexShapes,
+																const b3GpuChildShape* cpuChildShapes,
+																const b3AlignedObjectArray<b3Aabb>& hostAabbsWorldSpace,
+																const b3AlignedObjectArray<b3Aabb>& hostAabbsLocalSpace,
+
+																const b3AlignedObjectArray<b3Vector3>& convexVertices,
+																const b3AlignedObjectArray<b3Vector3>& hostUniqueEdges,
+																const b3AlignedObjectArray<int>& convexIndices,
+																const b3AlignedObjectArray<b3GpuFace>& faces,
+																
+																b3Contact4* globalContactsOut,
+																int& nGlobalContactsOut,
+																int maxContactCapacity)
+{
+
+	int shapeTypeB = collidables[collidableIndexB].m_shapeType;
+	b3Assert(shapeTypeB == SHAPE_COMPOUND_OF_CONVEX_HULLS);
+
+	b3AlignedObjectArray<b3Int4> cpuCompoundPairsOut;
+	int numCompoundPairsOut=0;
+	int maxNumCompoundPairsCapacity = 1024;
+	cpuCompoundPairsOut.resize(maxNumCompoundPairsCapacity);
+
+
+	// work-in-progress
+	findCompoundPairsKernel( 
+							pairIndex,
+							bodyIndexA,bodyIndexB,
+							collidableIndexA,collidableIndexB,
+							rigidBodies, 
+							collidables,
+							convexShapes, 
+							convexVertices,
+							hostAabbsWorldSpace,
+							hostAabbsLocalSpace,
+							cpuChildShapes,
+							&cpuCompoundPairsOut[0],
+							&numCompoundPairsOut,
+							maxNumCompoundPairsCapacity	);
+
+	b3AlignedObjectArray<b3Float4> cpuCompoundSepNormalsOut;
+	b3AlignedObjectArray<int> cpuHasCompoundSepNormalsOut;
+	cpuCompoundSepNormalsOut.resize(numCompoundPairsOut);
+	cpuHasCompoundSepNormalsOut.resize(numCompoundPairsOut);
+
+	for (int i=0;i<numCompoundPairsOut;i++)
+	{
+
+		processCompoundPairsKernel(&cpuCompoundPairsOut[0],rigidBodies,collidables,convexShapes,convexVertices,hostUniqueEdges,faces,convexIndices,0,cpuChildShapes,
+			cpuCompoundSepNormalsOut,cpuHasCompoundSepNormalsOut,numCompoundPairsOut,i);
+	}
+
+	for (int i=0;i<numCompoundPairsOut;i++)
+	{
+		clipCompoundsHullHullKernel(&cpuCompoundPairsOut[0],rigidBodies,collidables,convexShapes,convexVertices,hostUniqueEdges,faces,convexIndices,cpuChildShapes,
+			cpuCompoundSepNormalsOut,cpuHasCompoundSepNormalsOut,globalContactsOut,&nGlobalContactsOut,numCompoundPairsOut,maxContactCapacity,i);
+	}
+		/*
+		int childColIndexA = gpuChildShapes[childShapeIndexA].m_shapeIndex;
+
+					float4 posA = rigidBodies[bodyIndexA].m_pos;
+					b3Quat ornA = rigidBodies[bodyIndexA].m_quat;
+					float4 childPosA = gpuChildShapes[childShapeIndexA].m_childPosition;
+					b3Quat childOrnA = gpuChildShapes[childShapeIndexA].m_childOrientation;
+					float4 newPosA = b3QuatRotate(ornA,childPosA)+posA;
+					b3Quat newOrnA = b3QuatMul(ornA,childOrnA);
+
+					int shapeIndexA = collidables[childColIndexA].m_shapeIndex;
+
+
+			bool foundSepAxis = findSeparatingAxis(hullA,hullB,
+							posA,
+							ornA,
+							posB,
+							ornB,
+
+							convexVertices,uniqueEdges,faces,convexIndices,
+							convexVertices,uniqueEdges,faces,convexIndices,
+							
+							sepNormalWorldSpace
+							);
+							*/
+
+	
+	/*
+	if (foundSepAxis)
+	{
+		
+		
+		contactIndex = clipHullHullSingle(
+			bodyIndexA, bodyIndexB,
+						   posA,ornA,
+						   posB,ornB,
+			collidableIndexA, collidableIndexB,
+			&rigidBodies, 
+			&globalContactsOut,
+			nGlobalContactsOut,
+			
+			convexShapes,
+			convexShapes,
+	
+			convexVertices, 
+			uniqueEdges, 
+			faces,
+			convexIndices,
+	
+			convexVertices,
+			uniqueEdges,
+			faces,
+			convexIndices,
+
+			collidables,
+			collidables,
+			sepNormalWorldSpace,
+			maxContactCapacity);
+			
+	}
+	*/
+
+//	return contactIndex;
+	
+	/*
+
+	int numChildrenB = collidables[collidableIndexB].m_numChildShapes;
+	for (int c=0;c<numChildrenB;c++)
+	{
+		int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+c;
+		int childColIndexB = cpuChildShapes[childShapeIndexB].m_shapeIndex;
+
+		float4 rootPosB = rigidBodies[bodyIndexB].m_pos;
+		b3Quaternion rootOrnB = rigidBodies[bodyIndexB].m_quat;
+		b3Vector3 childPosB = cpuChildShapes[childShapeIndexB].m_childPosition;
+		b3Quaternion childOrnB = cpuChildShapes[childShapeIndexB].m_childOrientation;
+		float4  posB = b3QuatRotate(rootOrnB,childPosB)+rootPosB;
+		b3Quaternion ornB = b3QuatMul(rootOrnB,childOrnB);//b3QuatMul(ornB,childOrnB);
+
+		int shapeIndexB = collidables[childColIndexB].m_shapeIndex;
+
+		const b3ConvexPolyhedronCL* hullB = &convexShapes[shapeIndexB];
+
+	}
+	*/
+
+}
 
 void computeContactPlaneCompound(int pairIndex,
 																int bodyIndexA, int bodyIndexB, 
@@ -865,727 +2287,8 @@ void	computeContactSphereConvex(int pairIndex,
 	
 }
 
-#define MAX_VERTS 1024
 
 
-inline void project(const b3ConvexPolyhedronCL& hull,  const float4& pos, const b3Quaternion& orn, const float4& dir, const b3AlignedObjectArray<b3Vector3>& vertices, b3Scalar& min, b3Scalar& max)
-{
-	min = FLT_MAX;
-	max = -FLT_MAX;
-	int numVerts = hull.m_numVertices;
-
-	const float4 localDir = b3QuatRotate(orn.inverse(),dir);
-
-	b3Scalar offset = dot3F4(pos,dir);
-
-	for(int i=0;i<numVerts;i++)
-	{
-		//b3Vector3 pt = trans * vertices[m_vertexOffset+i];
-		//b3Scalar dp = pt.dot(dir);
-		b3Vector3 vertex = vertices[hull.m_vertexOffset+i];
-		b3Scalar dp = dot3F4((float4&)vertices[hull.m_vertexOffset+i],localDir);
-		//b3Assert(dp==dpL);
-		if(dp < min)	min = dp;
-		if(dp > max)	max = dp;
-	}
-	if(min>max)
-	{
-		b3Scalar tmp = min;
-		min = max;
-		max = tmp;
-	}
-	min += offset;
-	max += offset;
-}
-
-
-static bool TestSepAxis(const b3ConvexPolyhedronCL& hullA, const b3ConvexPolyhedronCL& hullB, 
-	const float4& posA,const b3Quaternion& ornA,
-	const float4& posB,const b3Quaternion& ornB,
-	const float4& sep_axis, const b3AlignedObjectArray<b3Vector3>& verticesA,const b3AlignedObjectArray<b3Vector3>& verticesB,b3Scalar& depth)
-{
-	b3Scalar Min0,Max0;
-	b3Scalar Min1,Max1;
-	project(hullA,posA,ornA,sep_axis,verticesA, Min0, Max0);
-	project(hullB,posB,ornB, sep_axis,verticesB, Min1, Max1);
-
-	if(Max0<Min1 || Max1<Min0)
-		return false;
-
-	b3Scalar d0 = Max0 - Min1;
-	assert(d0>=0.0f);
-	b3Scalar d1 = Max1 - Min0;
-	assert(d1>=0.0f);
-	depth = d0<d1 ? d0:d1;
-	return true;
-}
-
-inline bool IsAlmostZero(const b3Vector3& v)
-{
-	if(fabsf(v.x)>1e-6 || fabsf(v.y)>1e-6 || fabsf(v.z)>1e-6)	return false;
-	return true;
-}
-
-
-static bool findSeparatingAxis(	const b3ConvexPolyhedronCL& hullA, const b3ConvexPolyhedronCL& hullB, 
-	const float4& posA1,
-	const b3Quaternion& ornA,
-	const float4& posB1,
-	const b3Quaternion& ornB,
-	const b3AlignedObjectArray<b3Vector3>& verticesA,
-	const b3AlignedObjectArray<b3Vector3>& uniqueEdgesA, 
-	const b3AlignedObjectArray<b3GpuFace>& facesA,
-	const b3AlignedObjectArray<int>& indicesA,
-	const b3AlignedObjectArray<b3Vector3>& verticesB, 
-	const b3AlignedObjectArray<b3Vector3>& uniqueEdgesB, 
-	const b3AlignedObjectArray<b3GpuFace>& facesB,
-	const b3AlignedObjectArray<int>& indicesB,
-
-	b3Vector3& sep)
-{
-	B3_PROFILE("findSeparatingAxis");
-
-	b3g_actualSATPairTests++;
-	float4 posA = posA1;
-	posA.w = 0.f;
-	float4 posB = posB1;
-	posB.w = 0.f;
-//#ifdef TEST_INTERNAL_OBJECTS
-	float4 c0local = (float4&)hullA.m_localCenter;
-	float4 c0 = transform(&c0local, &posA, &ornA);
-	float4 c1local = (float4&)hullB.m_localCenter;
-	float4 c1 = transform(&c1local,&posB,&ornB);
-	const float4 deltaC2 = c0 - c1;
-//#endif
-
-	b3Scalar dmin = FLT_MAX;
-	int curPlaneTests=0;
-
-	int numFacesA = hullA.m_numFaces;
-	// Test normals from hullA
-	for(int i=0;i<numFacesA;i++)
-	{
-		const float4& normal = (float4&)facesA[hullA.m_faceOffset+i].m_plane;
-		float4 faceANormalWS = b3QuatRotate(ornA,normal);
-
-		if (dot3F4(deltaC2,faceANormalWS)<0)
-			faceANormalWS*=-1.f;
-
-		curPlaneTests++;
-#ifdef TEST_INTERNAL_OBJECTS
-		gExpectedNbTests++;
-		if(gUseInternalObject && !TestInternalObjects(transA,transB, DeltaC2, faceANormalWS, hullA, hullB, dmin))
-			continue;
-		gActualNbTests++;
-#endif
-
-		
-		b3Scalar d;
-		if(!TestSepAxis( hullA, hullB, posA,ornA,posB,ornB,faceANormalWS, verticesA, verticesB,d))
-			return false;
-
-		if(d<dmin)
-		{
-			dmin = d;
-			sep = (b3Vector3&)faceANormalWS;
-		}
-	}
-
-	int numFacesB = hullB.m_numFaces;
-	// Test normals from hullB
-	for(int i=0;i<numFacesB;i++)
-	{
-		float4 normal = (float4&)facesB[hullB.m_faceOffset+i].m_plane;
-		float4 WorldNormal = b3QuatRotate(ornB, normal);
-
-		if (dot3F4(deltaC2,WorldNormal)<0)
-		{
-			WorldNormal*=-1.f;
-		}
-		curPlaneTests++;
-#ifdef TEST_INTERNAL_OBJECTS
-		gExpectedNbTests++;
-		if(gUseInternalObject && !TestInternalObjects(transA,transB,DeltaC2, WorldNormal, hullA, hullB, dmin))
-			continue;
-		gActualNbTests++;
-#endif
-
-		b3Scalar d;
-		if(!TestSepAxis(hullA, hullB,posA,ornA,posB,ornB,WorldNormal,verticesA,verticesB,d))
-			return false;
-
-		if(d<dmin)
-		{
-			dmin = d;
-			sep = (b3Vector3&)WorldNormal;
-		}
-	}
-
-	b3Vector3 edgeAstart,edgeAend,edgeBstart,edgeBend;
-
-	int curEdgeEdge = 0;
-	// Test edges
-	for(int e0=0;e0<hullA.m_numUniqueEdges;e0++)
-	{
-		const float4& edge0 = (float4&) uniqueEdgesA[hullA.m_uniqueEdgesOffset+e0];
-		float4 edge0World = b3QuatRotate(ornA,(float4&)edge0);
-
-		for(int e1=0;e1<hullB.m_numUniqueEdges;e1++)
-		{
-			const b3Vector3 edge1 = uniqueEdgesB[hullB.m_uniqueEdgesOffset+e1];
-			float4 edge1World = b3QuatRotate(ornB,(float4&)edge1);
-
-
-			float4 crossje = cross3(edge0World,edge1World);
-
-			curEdgeEdge++;
-			if(!IsAlmostZero((b3Vector3&)crossje))
-			{
-				crossje = normalize3(crossje);
-				if (dot3F4(deltaC2,crossje)<0)
-					crossje*=-1.f;
-
-
-#ifdef TEST_INTERNAL_OBJECTS
-				gExpectedNbTests++;
-				if(gUseInternalObject && !TestInternalObjects(transA,transB,DeltaC2, Cross, hullA, hullB, dmin))
-					continue;
-				gActualNbTests++;
-#endif
-
-				b3Scalar dist;
-				if(!TestSepAxis( hullA, hullB, posA,ornA,posB,ornB,crossje, verticesA,verticesB,dist))
-					return false;
-
-				if(dist<dmin)
-				{
-					dmin = dist;
-					sep = (b3Vector3&)crossje;
-				}
-			}
-		}
-
-	}
-
-	
-	if((dot3F4(-deltaC2,(float4&)sep))>0.0f)
-		sep = -sep;
-
-	return true;
-}
-
-
-__inline float4 lerp3(const float4& a,const float4& b, float  t)
-{
-	return make_float4(	a.x + (b.x - a.x) * t,
-						a.y + (b.y - a.y) * t,
-						a.z + (b.z - a.z) * t,
-						0.f);
-}
-
-
-// Clips a face to the back of a plane, return the number of vertices out, stored in ppVtxOut
-int clipFace(const float4* pVtxIn, int numVertsIn, float4& planeNormalWS,float planeEqWS, float4* ppVtxOut)
-{
-	
-	int ve;
-	float ds, de;
-	int numVertsOut = 0;
-	if (numVertsIn < 2)
-		return 0;
-
-	float4 firstVertex=pVtxIn[numVertsIn-1];
-	float4 endVertex = pVtxIn[0];
-	
-	ds = dot3F4(planeNormalWS,firstVertex)+planeEqWS;
-
-	for (ve = 0; ve < numVertsIn; ve++)
-	{
-		endVertex=pVtxIn[ve];
-
-		de = dot3F4(planeNormalWS,endVertex)+planeEqWS;
-
-		if (ds<0)
-		{
-			if (de<0)
-			{
-				// Start < 0, end < 0, so output endVertex
-				ppVtxOut[numVertsOut++] = endVertex;
-			}
-			else
-			{
-				// Start < 0, end >= 0, so output intersection
-				ppVtxOut[numVertsOut++] = lerp3(firstVertex, endVertex,(ds * 1.f/(ds - de)) );
-			}
-		}
-		else
-		{
-			if (de<0)
-			{
-				// Start >= 0, end < 0 so output intersection and end
-				ppVtxOut[numVertsOut++] = lerp3(firstVertex, endVertex,(ds * 1.f/(ds - de)) );
-				ppVtxOut[numVertsOut++] = endVertex;
-			}
-		}
-		firstVertex = endVertex;
-		ds = de;
-	}
-	return numVertsOut;
-}
-
-
-int clipFaceAgainstHull(const float4& separatingNormal, const b3ConvexPolyhedronCL* hullA,  
-	const float4& posA, const b3Quaternion& ornA, float4* worldVertsB1, int numWorldVertsB1,
-	float4* worldVertsB2, int capacityWorldVertsB2,
-	const float minDist, float maxDist,
-	const float4* verticesA,	const b3GpuFace* facesA,	const int* indicesA,
-	//const float4* verticesB,	const b3GpuFace* facesB,	const int* indicesB,
-	float4* contactsOut,
-	int contactCapacity)
-{
-	int numContactsOut = 0;
-
-	float4* pVtxIn = worldVertsB1;
-	float4* pVtxOut = worldVertsB2;
-	
-	int numVertsIn = numWorldVertsB1;
-	int numVertsOut = 0;
-
-	int closestFaceA=-1;
-	{
-		float dmin = FLT_MAX;
-		for(int face=0;face<hullA->m_numFaces;face++)
-		{
-			const float4 Normal = make_float4(
-				facesA[hullA->m_faceOffset+face].m_plane.x, 
-				facesA[hullA->m_faceOffset+face].m_plane.y, 
-				facesA[hullA->m_faceOffset+face].m_plane.z,0.f);
-			const float4 faceANormalWS = b3QuatRotate(ornA,Normal);
-		
-			float d = dot3F4(faceANormalWS,separatingNormal);
-			if (d < dmin)
-			{
-				dmin = d;
-				closestFaceA = face;
-			}
-		}
-	}
-	if (closestFaceA<0)
-		return numContactsOut;
-
-	b3GpuFace polyA = facesA[hullA->m_faceOffset+closestFaceA];
-
-	// clip polygon to back of planes of all faces of hull A that are adjacent to witness face
-	int numContacts = numWorldVertsB1;
-	int numVerticesA = polyA.m_numIndices;
-	for(int e0=0;e0<numVerticesA;e0++)
-	{
-		const float4 a = verticesA[hullA->m_vertexOffset+indicesA[polyA.m_indexOffset+e0]];
-		const float4 b = verticesA[hullA->m_vertexOffset+indicesA[polyA.m_indexOffset+((e0+1)%numVerticesA)]];
-		const float4 edge0 = a - b;
-		const float4 WorldEdge0 = b3QuatRotate(ornA,edge0);
-		float4 planeNormalA = make_float4(polyA.m_plane.x,polyA.m_plane.y,polyA.m_plane.z,0.f);
-		float4 worldPlaneAnormal1 = b3QuatRotate(ornA,planeNormalA);
-
-		float4 planeNormalWS1 = -cross3(WorldEdge0,worldPlaneAnormal1);
-		float4 worldA1 = transform(&a,&posA,&ornA);
-		float planeEqWS1 = -dot3F4(worldA1,planeNormalWS1);
-		
-		float4 planeNormalWS = planeNormalWS1;
-		float planeEqWS=planeEqWS1;
-		
-		//clip face
-		//clipFace(*pVtxIn, *pVtxOut,planeNormalWS,planeEqWS);
-		numVertsOut = clipFace(pVtxIn, numVertsIn, planeNormalWS,planeEqWS, pVtxOut);
-
-		//btSwap(pVtxIn,pVtxOut);
-		float4* tmp = pVtxOut;
-		pVtxOut = pVtxIn;
-		pVtxIn = tmp;
-		numVertsIn = numVertsOut;
-		numVertsOut = 0;
-	}
-
-	
-	// only keep points that are behind the witness face
-	{
-		float4 localPlaneNormal  = make_float4(polyA.m_plane.x,polyA.m_plane.y,polyA.m_plane.z,0.f);
-		float localPlaneEq = polyA.m_plane.w;
-		float4 planeNormalWS = b3QuatRotate(ornA,localPlaneNormal);
-		float planeEqWS=localPlaneEq-dot3F4(planeNormalWS,posA);
-		for (int i=0;i<numVertsIn;i++)
-		{
-			float depth = dot3F4(planeNormalWS,pVtxIn[i])+planeEqWS;
-			if (depth <=minDist)
-			{
-				depth = minDist;
-			}
-			if (numContactsOut<contactCapacity)
-			{
-				if (depth <=maxDist)
-				{
-					float4 pointInWorld = pVtxIn[i];
-					//resultOut.addContactPoint(separatingNormal,point,depth);
-					contactsOut[numContactsOut++] = make_float4(pointInWorld.x,pointInWorld.y,pointInWorld.z,depth);
-					//printf("depth=%f\n",depth);
-				}
-			} else
-			{
-				b3Error("exceeding contact capacity (%d,%df)\n", numContactsOut,contactCapacity);
-			}
-		}
-	}
-
-	return numContactsOut;
-}
-
-
-
-static int	clipHullAgainstHull(const float4& separatingNormal, 
-	const b3ConvexPolyhedronCL& hullA, const b3ConvexPolyhedronCL& hullB, 
-	const float4& posA, const b3Quaternion& ornA,const float4& posB, const b3Quaternion& ornB, 
-	float4* worldVertsB1, float4* worldVertsB2, int capacityWorldVerts,
-	const float minDist, float maxDist,
-	const float4* verticesA,	const b3GpuFace* facesA,	const int* indicesA,
-	const float4* verticesB,	const b3GpuFace* facesB,	const int* indicesB,
-
-	float4*	contactsOut,
-	int contactCapacity)
-{
-	int numContactsOut = 0;
-	int numWorldVertsB1= 0;
-	
-	B3_PROFILE("clipHullAgainstHull");
-
-	float curMaxDist=maxDist;
-	int closestFaceB=-1;
-	float dmax = -FLT_MAX;
-
-	{
-		//B3_PROFILE("closestFaceB");
-		if (hullB.m_numFaces!=1)
-		{
-			//printf("wtf\n");
-		}
-		static bool once = true;
-		//printf("separatingNormal=%f,%f,%f\n",separatingNormal.x,separatingNormal.y,separatingNormal.z);
-		
-		for(int face=0;face<hullB.m_numFaces;face++)
-		{
-#ifdef BT_DEBUG_SAT_FACE
-			if (once)
-				printf("face %d\n",face);
-			const b3GpuFace* faceB = &facesB[hullB.m_faceOffset+face];
-			if (once)
-			{
-				for (int i=0;i<faceB->m_numIndices;i++)
-				{
-					float4 vert = verticesB[hullB.m_vertexOffset+indicesB[faceB->m_indexOffset+i]];
-					printf("vert[%d] = %f,%f,%f\n",i,vert.x,vert.y,vert.z);
-				}
-			}
-#endif //BT_DEBUG_SAT_FACE
-			//if (facesB[hullB.m_faceOffset+face].m_numIndices>2)
-			{
-				const float4 Normal = make_float4(facesB[hullB.m_faceOffset+face].m_plane.x, 
-					facesB[hullB.m_faceOffset+face].m_plane.y, facesB[hullB.m_faceOffset+face].m_plane.z,0.f);
-				const float4 WorldNormal = b3QuatRotate(ornB, Normal);
-#ifdef BT_DEBUG_SAT_FACE
-				if (once)
-					printf("faceNormal = %f,%f,%f\n",Normal.x,Normal.y,Normal.z);
-#endif
-				float d = dot3F4(WorldNormal,separatingNormal);
-				if (d > dmax)
-				{
-					dmax = d;
-					closestFaceB = face;
-				}
-			}
-		}
-		once = false;
-	}
-
-	
-	b3Assert(closestFaceB>=0);
-	{
-		//B3_PROFILE("worldVertsB1");
-		const b3GpuFace& polyB = facesB[hullB.m_faceOffset+closestFaceB];
-		const int numVertices = polyB.m_numIndices;
-		for(int e0=0;e0<numVertices;e0++)
-		{
-			const float4& b = verticesB[hullB.m_vertexOffset+indicesB[polyB.m_indexOffset+e0]];
-			worldVertsB1[numWorldVertsB1++] = transform(&b,&posB,&ornB);
-		}
-	}
-
-	if (closestFaceB>=0)
-	{
-		//B3_PROFILE("clipFaceAgainstHull");
-		numContactsOut = clipFaceAgainstHull((float4&)separatingNormal, &hullA, 
-				posA,ornA,
-				worldVertsB1,numWorldVertsB1,worldVertsB2,capacityWorldVerts, minDist, maxDist,
-				verticesA,				facesA,				indicesA,
-				contactsOut,contactCapacity);
-	}
-
-	return numContactsOut;
-}
-
-
-
-
-
-
-#define PARALLEL_SUM(v, n) for(int j=1; j<n; j++) v[0] += v[j];
-#define PARALLEL_DO(execution, n) for(int ie=0; ie<n; ie++){execution;}
-#define REDUCE_MAX(v, n) {int i=0;\
-for(int offset=0; offset<n; offset++) v[i] = (v[i].y > v[i+offset].y)? v[i]: v[i+offset]; }
-#define REDUCE_MIN(v, n) {int i=0;\
-for(int offset=0; offset<n; offset++) v[i] = (v[i].y < v[i+offset].y)? v[i]: v[i+offset]; }
-
-int extractManifold(const float4* p, int nPoints, const float4& nearNormal, b3Int4* contactIdx)
-{
-	if( nPoints == 0 )
-        return 0;
-    
-    if (nPoints <=4)
-        return nPoints;
-    
-    
-    if (nPoints >64)
-        nPoints = 64;
-    
-	float4 center = make_float4(0,0,0,0);
-	{
-		
-		for (int i=0;i<nPoints;i++)
-			center += p[i];
-		center /= (float)nPoints;
-	}
-    
-	
-    
-	//	sample 4 directions
-    
-    float4 aVector = p[0] - center;
-    float4 u = cross3( nearNormal, aVector );
-    float4 v = cross3( nearNormal, u );
-    u = normalize3( u );
-    v = normalize3( v );
-    
-    
-    //keep point with deepest penetration
-    float minW= FLT_MAX;
-    
-    int minIndex=-1;
-    
-    float4 maxDots;
-    maxDots.x = FLT_MIN;
-    maxDots.y = FLT_MIN;
-    maxDots.z = FLT_MIN;
-    maxDots.w = FLT_MIN;
-    
-    //	idx, distance
-    for(int ie = 0; ie<nPoints; ie++ )
-    {
-        if (p[ie].w<minW)
-        {
-            minW = p[ie].w;
-            minIndex=ie;
-        }
-        float f;
-        float4 r = p[ie]-center;
-        f = dot3F4( u, r );
-        if (f<maxDots.x)
-        {
-            maxDots.x = f;
-            contactIdx[0].x = ie;
-        }
-        
-        f = dot3F4( -u, r );
-        if (f<maxDots.y)
-        {
-            maxDots.y = f;
-            contactIdx[0].y = ie;
-        }
-        
-        
-        f = dot3F4( v, r );
-        if (f<maxDots.z)
-        {
-            maxDots.z = f;
-            contactIdx[0].z = ie;
-        }
-        
-        f = dot3F4( -v, r );
-        if (f<maxDots.w)
-        {
-            maxDots.w = f;
-            contactIdx[0].w = ie;
-        }
-        
-    }
-    
-    if (contactIdx[0].x != minIndex && contactIdx[0].y != minIndex && contactIdx[0].z != minIndex && contactIdx[0].w != minIndex)
-    {
-        //replace the first contact with minimum (todo: replace contact with least penetration)
-        contactIdx[0].x = minIndex;
-    }
-    
-    return 4;
-    
-}
-
-
-
-
-int clipHullHullSingle(
-			int bodyIndexA, int bodyIndexB,
-										 const float4& posA,
-										 const b3Quaternion& ornA,
-										 const float4& posB,
-										 const b3Quaternion& ornB,
-
-			int collidableIndexA, int collidableIndexB,
-
-			const b3AlignedObjectArray<b3RigidBodyCL>* bodyBuf, 
-			b3AlignedObjectArray<b3Contact4>* globalContactOut, 
-			int& nContacts,
-			
-			const b3AlignedObjectArray<b3ConvexPolyhedronCL>& hostConvexDataA,
-			const b3AlignedObjectArray<b3ConvexPolyhedronCL>& hostConvexDataB,
-	
-			const b3AlignedObjectArray<b3Vector3>& verticesA, 
-			const b3AlignedObjectArray<b3Vector3>& uniqueEdgesA, 
-			const b3AlignedObjectArray<b3GpuFace>& facesA,
-			const b3AlignedObjectArray<int>& indicesA,
-	
-			const b3AlignedObjectArray<b3Vector3>& verticesB,
-			const b3AlignedObjectArray<b3Vector3>& uniqueEdgesB,
-			const b3AlignedObjectArray<b3GpuFace>& facesB,
-			const b3AlignedObjectArray<int>& indicesB,
-
-			const b3AlignedObjectArray<b3Collidable>& hostCollidablesA,
-			const b3AlignedObjectArray<b3Collidable>& hostCollidablesB,
-			const b3Vector3& sepNormalWorldSpace,
-			int maxContactCapacity			)
-{
-	int contactIndex = -1;
-	b3ConvexPolyhedronCL hullA, hullB;
-    
-    b3Collidable colA = hostCollidablesA[collidableIndexA];
-    hullA = hostConvexDataA[colA.m_shapeIndex];
-    //printf("numvertsA = %d\n",hullA.m_numVertices);
-    
-    
-    b3Collidable colB = hostCollidablesB[collidableIndexB];
-    hullB = hostConvexDataB[colB.m_shapeIndex];
-    //printf("numvertsB = %d\n",hullB.m_numVertices);
-    
-	
-	float4 contactsOut[MAX_VERTS];
-	int localContactCapacity = MAX_VERTS;
-
-#ifdef _WIN32
-	b3Assert(_finite(bodyBuf->at(bodyIndexA).m_pos.x));
-	b3Assert(_finite(bodyBuf->at(bodyIndexB).m_pos.x));
-#endif
-	
-	
-	{
-		
-		float4 worldVertsB1[MAX_VERTS];
-		float4 worldVertsB2[MAX_VERTS];
-		int capacityWorldVerts = MAX_VERTS;
-
-		float4 hostNormal = make_float4(sepNormalWorldSpace.getX(),sepNormalWorldSpace.getY(),sepNormalWorldSpace.getZ(),0.f);
-		int shapeA = hostCollidablesA[collidableIndexA].m_shapeIndex;
-		int shapeB = hostCollidablesB[collidableIndexB].m_shapeIndex;
-
-		b3Scalar minDist = -1;
-		b3Scalar maxDist = 0.;
-
-		        
-
-		b3Transform trA,trB;
-		{
-		//B3_PROFILE("transform computation");
-		//trA.setIdentity();
-		trA.setOrigin(b3Vector3(posA.x,posA.y,posA.z));
-		trA.setRotation(b3Quaternion(ornA.x,ornA.y,ornA.z,ornA.w));
-				
-		//trB.setIdentity();
-		trB.setOrigin(b3Vector3(posB.x,posB.y,posB.z));
-		trB.setRotation(b3Quaternion(ornB.x,ornB.y,ornB.z,ornB.w));
-		}
-
-		b3Quaternion trAorn = trA.getRotation();
-        b3Quaternion trBorn = trB.getRotation();
-        
-		int numContactsOut = clipHullAgainstHull(hostNormal, 
-						hostConvexDataA.at(shapeA), 
-						hostConvexDataB.at(shapeB),
-								(float4&)trA.getOrigin(), (b3Quaternion&)trAorn,
-								(float4&)trB.getOrigin(), (b3Quaternion&)trBorn,
-								worldVertsB1,worldVertsB2,capacityWorldVerts,
-								minDist, maxDist,
-								(float4*)&verticesA[0],	&facesA[0],&indicesA[0],
-								(float4*)&verticesB[0],	&facesB[0],&indicesB[0],
-								
-								contactsOut,localContactCapacity);
-
-		if (numContactsOut>0)
-		{
-			B3_PROFILE("overlap");
-
-			float4 normalOnSurfaceB = (float4&)hostNormal;
-			float4 centerOut;
-			
-			b3Int4 contactIdx;
-			contactIdx.x = 0;
-			contactIdx.y = 1;
-			contactIdx.z = 2;
-			contactIdx.w = 3;
-			
-			int numPoints = 0;
-					
-			{
-				B3_PROFILE("extractManifold");
-				numPoints = extractManifold(contactsOut, numContactsOut, normalOnSurfaceB,  &contactIdx);
-			}
-					
-			b3Assert(numPoints);
-					
-			if (nContacts<maxContactCapacity)
-			{
-				contactIndex = nContacts;
-				globalContactOut->expand();
-				b3Contact4& contact = globalContactOut->at(nContacts);
-				contact.m_batchIdx = 0;//i;
-				contact.m_bodyAPtrAndSignBit = (bodyBuf->at(bodyIndexA).m_invMass==0)? -bodyIndexA:bodyIndexA;
-				contact.m_bodyBPtrAndSignBit = (bodyBuf->at(bodyIndexB).m_invMass==0)? -bodyIndexB:bodyIndexB;
-
-				contact.m_frictionCoeffCmp = 45874;
-				contact.m_restituitionCoeffCmp = 0;
-					
-				float distance = 0.f;
-				for (int p=0;p<numPoints;p++)
-				{
-					contact.m_worldPosB[p] = contactsOut[contactIdx.s[p]];//check if it is actually on B
-					contact.m_worldNormalOnB = normalOnSurfaceB; 
-				}
-				//printf("bodyIndexA %d,bodyIndexB %d,normal=%f,%f,%f numPoints %d\n",bodyIndexA,bodyIndexB,normalOnSurfaceB.x,normalOnSurfaceB.y,normalOnSurfaceB.z,numPoints);
-				contact.m_worldNormalOnB.w = (b3Scalar)numPoints;
-				nContacts++;
-			} else
-			{
-				b3Error("Error: exceeding contact capacity (%d/%d)\n", nContacts,maxContactCapacity);
-			}
-		}
-	}
-	return contactIndex;
-}
-
 #include "b3GjkPairDetector.h"
 #include "b3GjkEpa.h"
 #include "b3VoronoiSimplexSolver.h"
@@ -1832,7 +2535,9 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
 			const b3OpenCLArray<b3Collidable>& gpuCollidables,
 			const b3OpenCLArray<b3GpuChildShape>& gpuChildShapes,
 
-			const b3OpenCLArray<b3YetAnotherAabb>& clAabbsWS,
+			const b3OpenCLArray<b3Aabb>& clAabbsWorldSpace,
+			const b3OpenCLArray<b3Aabb>& clAabbslocalSpace,
+
             b3OpenCLArray<b3Vector3>& worldVertsB1GPU,
             b3OpenCLArray<b3Int4>& clippingFacesOutGPU,
             b3OpenCLArray<b3Vector3>& worldNormalsAGPU,
@@ -1855,8 +2560,12 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
 
 
 #ifdef CHECK_ON_HOST
-	b3AlignedObjectArray<b3YetAnotherAabb> hostAabbs;
-	clAabbsWS.copyToHost(hostAabbs);
+	b3AlignedObjectArray<b3Aabb> hostAabbsWorldSpace;
+	clAabbsWorldSpace.copyToHost(hostAabbsWorldSpace);
+
+	b3AlignedObjectArray<b3Aabb> hostAabbsLocalSpace;
+	clAabbslocalSpace.copyToHost(hostAabbsLocalSpace);
+
 	b3AlignedObjectArray<b3Int4> hostPairs;
 	pairs->copyToHost(hostPairs);
 
@@ -1940,6 +2649,15 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
 			&hostCollidables[0],&hostConvexData[0],&hostVertices[0],&hostIndices[0],&hostFaces[0],&hostContacts[0],nContacts,maxContactCapacity);
 //			printf("plane-convex\n");
 			
+		}
+
+			if (hostCollidables[collidableIndexA].m_shapeType == SHAPE_COMPOUND_OF_CONVEX_HULLS &&
+			hostCollidables[collidableIndexB].m_shapeType == SHAPE_COMPOUND_OF_CONVEX_HULLS)
+		{
+			computeContactCompoundCompound(i,bodyIndexB,bodyIndexA,collidableIndexB,collidableIndexA,&hostBodyBuf[0],
+			&hostCollidables[0],&hostConvexData[0],&cpuChildShapes[0], hostAabbsWorldSpace,hostAabbsLocalSpace,hostVertices,hostUniqueEdges,hostIndices,hostFaces,&hostContacts[0],nContacts,maxContactCapacity);
+//			printf("convex-plane\n");
+			
 		}
 
 
@@ -2079,7 +2797,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
 					b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
 					b3BufferInfoCL( gpuFaces.getBufferCL(),true),
 					b3BufferInfoCL( gpuIndices.getBufferCL(),true),
-					b3BufferInfoCL( clAabbsWS.getBufferCL(),true),
+					b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
 					b3BufferInfoCL( m_sepNormals.getBufferCL()),
 					b3BufferInfoCL( m_hasSeparatingNormals.getBufferCL())
 				};
@@ -2111,7 +2829,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
 						launcher.setBuffer( pairs->getBufferCL());
 						launcher.setBuffer(  bodyBuf->getBufferCL());
 						launcher.setBuffer( gpuCollidables.getBufferCL());
-						launcher.setBuffer( clAabbsWS.getBufferCL());
+						launcher.setBuffer( clAabbsWorldSpace.getBufferCL());
 						launcher.setBuffer( triangleConvexPairsOut.getBufferCL());
 						launcher.setBuffer( m_numConcavePairsOut.getBufferCL());
 						launcher.setBuffer( subTreesGPU->getBufferCL());
@@ -2147,7 +2865,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
 								b3BufferInfoCL( gpuFaces.getBufferCL(),true),
 								b3BufferInfoCL( gpuIndices.getBufferCL(),true),
 								b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
-								b3BufferInfoCL( clAabbsWS.getBufferCL(),true),
+								b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
 								b3BufferInfoCL( m_concaveSepNormals.getBufferCL())
 							};
 
@@ -2186,7 +2904,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
 					b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
 					b3BufferInfoCL( gpuFaces.getBufferCL(),true),
 					b3BufferInfoCL( gpuIndices.getBufferCL(),true),
-					b3BufferInfoCL( clAabbsWS.getBufferCL(),true),
+					b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
 					b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
 					b3BufferInfoCL( m_gpuCompoundPairs.getBufferCL()),
 					b3BufferInfoCL( m_numCompoundPairsOut.getBufferCL())
@@ -2230,7 +2948,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
 					b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
 					b3BufferInfoCL( gpuFaces.getBufferCL(),true),
 					b3BufferInfoCL( gpuIndices.getBufferCL(),true),
-					b3BufferInfoCL( clAabbsWS.getBufferCL(),true),
+					b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
 					b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
 					b3BufferInfoCL( contactOut->getBufferCL()),
 					b3BufferInfoCL( m_totalContactsOut.getBufferCL())	
@@ -2269,7 +2987,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
 					b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
 					b3BufferInfoCL( gpuFaces.getBufferCL(),true),
 					b3BufferInfoCL( gpuIndices.getBufferCL(),true),
-					b3BufferInfoCL( clAabbsWS.getBufferCL(),true),
+					b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
 					b3BufferInfoCL( gpuChildShapes.getBufferCL(),true),
 					b3BufferInfoCL( m_gpuCompoundSepNormals.getBufferCL()),
 					b3BufferInfoCL( m_gpuHasCompoundSepNormals.getBufferCL())
@@ -2316,7 +3034,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
 				b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
 				b3BufferInfoCL( gpuFaces.getBufferCL(),true),
 				b3BufferInfoCL( gpuIndices.getBufferCL(),true),
-				b3BufferInfoCL( clAabbsWS.getBufferCL(),true),
+				b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
 				b3BufferInfoCL( contactOut->getBufferCL()),
 				b3BufferInfoCL( m_totalContactsOut.getBufferCL())
 			};
diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.h b/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.h
index a78063766..f28027768 100644
--- a/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.h
+++ b/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.h
@@ -7,30 +7,18 @@
 #include "Bullet3Common/b3AlignedObjectArray.h"
 #include "b3ConvexUtility.h"
 #include "b3ConvexPolyhedronCL.h"
-#include "b3Collidable.h"
+#include "Bullet3Collision/NarrowPhaseCollision/shared/b3Collidable.h"
 #include "Bullet3Collision/NarrowPhaseCollision/b3Contact4.h"
-#include "Bullet3Common/b3Int2.h"
-#include "Bullet3Common/b3Int4.h"
+#include "Bullet3Common/shared/b3Int2.h"
+#include "Bullet3Common/shared/b3Int4.h"
 #include "b3OptimizedBvh.h"
 #include "b3BvhInfo.h"
+#include "Bullet3Collision/BroadPhaseCollision/shared/b3Aabb.h"
 
 //#include "../../dynamics/basic_demo/Stubs/ChNarrowPhase.h"
 
 
-struct b3YetAnotherAabb
-{
-	union
-	{
-		float m_min[4];
-		int m_minIndices[4];
-	};
-	union
-	{
-		float m_max[4];
-		//int m_signedMaxIndices[4];
-		//unsigned int m_unsignedMaxIndices[4];
-	};
-};
+
 
 struct GpuSatCollision
 {
@@ -89,7 +77,9 @@ struct GpuSatCollision
 			const b3OpenCLArray<b3Collidable>& gpuCollidables,
 			const b3OpenCLArray<b3GpuChildShape>& gpuChildShapes,
 
-			const b3OpenCLArray<b3YetAnotherAabb>& clAabbs,
+			const b3OpenCLArray<b3Aabb>& clAabbsWorldSpace,
+			const b3OpenCLArray<b3Aabb>& clAabbsLocalSpace,
+
            b3OpenCLArray<b3Vector3>& worldVertsB1GPU,
            b3OpenCLArray<b3Int4>& clippingFacesOutGPU,
            b3OpenCLArray<b3Vector3>& worldNormalsAGPU,
diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexPolyhedronCL.h b/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexPolyhedronCL.h
index 97c525b1d..b13dcd945 100644
--- a/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexPolyhedronCL.h
+++ b/src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexPolyhedronCL.h
@@ -2,34 +2,12 @@
 #define CONVEX_POLYHEDRON_CL
 
 #include "Bullet3Common/b3Transform.h"
+#include "Bullet3Collision/NarrowPhaseCollision/shared/b3ConvexPolyhedronData.h"
 
-struct b3GpuFace
+
+
+B3_ATTRIBUTE_ALIGNED16(struct) b3ConvexPolyhedronCL : public b3ConvexPolyhedronData
 {
-	b3Vector4 m_plane;
-	int m_indexOffset;
-	int m_numIndices;
-	int m_unusedPadding1;
-	int m_unusedPadding2;
-};
-
-B3_ATTRIBUTE_ALIGNED16(struct) b3ConvexPolyhedronCL
-{
-	b3Vector3		m_localCenter;
-	b3Vector3		m_extents;
-	b3Vector3		mC;
-	b3Vector3		mE;
-
-	b3Scalar		m_radius;
-	int	m_faceOffset;
-	int m_numFaces;
-	int	m_numVertices;
-
-	int m_vertexOffset;
-	int	m_uniqueEdgesOffset;
-	int	m_numUniqueEdges;
-	int m_unused;	
-	
-
 
 	inline void project(const b3Transform& trans, const b3Vector3& dir, const b3AlignedObjectArray<b3Vector3>& vertices, b3Scalar& min, b3Scalar& max) const
 	{
diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.h b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.h
index 7c7dd35ed..eb183b82f 100644
--- a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.h
+++ b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/primitiveContacts.h
@@ -10,9 +10,13 @@ static const char* primitiveContactsKernelsCL= \
 "{\n"
 "	int bla;\n"
 "};\n"
+"#ifdef __cplusplus\n"
+"#else\n"
+"#define b3AtomicInc atomic_inc\n"
+"#endif\n"
 "#endif\n"
 "#ifdef __cplusplus\n"
-"#else//bla\n"
+"#else\n"
 "	typedef float4	b3Float4;\n"
 "#endif \n"
 "#endif //B3_FLOAT4_H\n"
diff --git a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.h b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.h
index 8d7a6ec6d..21106a8b8 100644
--- a/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.h
+++ b/src/Bullet3OpenCL/NarrowphaseCollision/kernels/satClipHullContacts.h
@@ -42,9 +42,13 @@ static const char* satClipKernelsCL= \
 "{\n"
 "	int bla;\n"
 "};\n"
+"#ifdef __cplusplus\n"
+"#else\n"
+"#define b3AtomicInc atomic_inc\n"
+"#endif\n"
 "#endif\n"
 "#ifdef __cplusplus\n"
-"#else//bla\n"
+"#else\n"
 "	typedef float4	b3Float4;\n"
 "#endif \n"
 "#endif //B3_FLOAT4_H\n"
diff --git a/src/Bullet3OpenCL/ParallelPrimitives/b3FillCL.h b/src/Bullet3OpenCL/ParallelPrimitives/b3FillCL.h
index 3803d202c..1609676b9 100644
--- a/src/Bullet3OpenCL/ParallelPrimitives/b3FillCL.h
+++ b/src/Bullet3OpenCL/ParallelPrimitives/b3FillCL.h
@@ -4,8 +4,8 @@
 #include "b3OpenCLArray.h"
 #include "Bullet3Common/b3Scalar.h"
 
-#include "Bullet3Common/b3Int2.h"
-#include "Bullet3Common/b3Int4.h"
+#include "Bullet3Common/shared/b3Int2.h"
+#include "Bullet3Common/shared/b3Int4.h"
 
 
 class b3FillCL
diff --git a/src/Bullet3OpenCL/Raycast/b3GpuRaycast.cpp b/src/Bullet3OpenCL/Raycast/b3GpuRaycast.cpp
index 8f0b77b3f..4e110aeb6 100644
--- a/src/Bullet3OpenCL/Raycast/b3GpuRaycast.cpp
+++ b/src/Bullet3OpenCL/Raycast/b3GpuRaycast.cpp
@@ -1,6 +1,6 @@
 
 #include "b3GpuRaycast.h"
-#include "Bullet3OpenCL/NarrowphaseCollision/b3Collidable.h"
+#include "Bullet3Collision/NarrowPhaseCollision/shared/b3Collidable.h"
 #include "Bullet3Collision/NarrowPhaseCollision/b3RigidBodyCL.h"
 #include "Bullet3OpenCL/RigidBody/b3GpuNarrowPhaseInternalData.h"
 
diff --git a/src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.cpp b/src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.cpp
index e6dca3085..faa468086 100644
--- a/src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.cpp
+++ b/src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.cpp
@@ -698,7 +698,7 @@ const struct b3Collidable* b3GpuNarrowPhase::getCollidablesCpu() const
 }
 
 
-cl_mem	b3GpuNarrowPhase::getAabbBufferGpu()
+cl_mem	b3GpuNarrowPhase::getAabbLocalSpaceBufferGpu()
 {
 	return m_data->m_localShapeAABBGPU->getBufferCL();
 }
@@ -726,8 +726,11 @@ const b3Contact4* b3GpuNarrowPhase::getContactsCPU() const
 	return &m_data->m_pBufContactOutCPU->at(0);
 }
 
-void b3GpuNarrowPhase::computeContacts(cl_mem broadphasePairs, int numBroadphasePairs, cl_mem aabbsWS, int numObjects)
+void b3GpuNarrowPhase::computeContacts(cl_mem broadphasePairs, int numBroadphasePairs, cl_mem aabbsWorldSpace, int numObjects)
 {
+
+	cl_mem aabbsLocalSpace = m_data->m_localShapeAABBGPU->getBufferCL();
+
 	int nContactOut = 0;
 
 	//swap buffer
@@ -744,8 +747,11 @@ void b3GpuNarrowPhase::computeContacts(cl_mem broadphasePairs, int numBroadphase
 	
 
 
-	b3OpenCLArray<b3YetAnotherAabb> clAabbArray(this->m_context,this->m_queue);
-	clAabbArray.setFromOpenCLBuffer(aabbsWS,numObjects);
+	b3OpenCLArray<b3Aabb> clAabbArrayWorldSpace(this->m_context,this->m_queue);
+	clAabbArrayWorldSpace.setFromOpenCLBuffer(aabbsWorldSpace,numObjects);
+
+	b3OpenCLArray<b3Aabb> clAabbArrayLocalSpace(this->m_context,this->m_queue);
+	clAabbArrayLocalSpace.setFromOpenCLBuffer(aabbsLocalSpace,numObjects);
 
 	m_data->m_gpuSatCollision->computeConvexConvexContactsGPUSAT(
 		&broadphasePairsGPU, numBroadphasePairs,
@@ -762,7 +768,8 @@ void b3GpuNarrowPhase::computeContacts(cl_mem broadphasePairs, int numBroadphase
 		*m_data->m_convexIndicesGPU,
 		*m_data->m_collidablesGPU,
 		*m_data->m_gpuChildShapes,
-		clAabbArray,
+		clAabbArrayWorldSpace,
+		clAabbArrayLocalSpace,
 		*m_data->m_worldVertsB1GPU,
 		*m_data->m_clippingFacesOutGPU,
 		*m_data->m_worldNormalsAGPU,
diff --git a/src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.h b/src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.h
index a916bbe57..3f7cd2e7e 100644
--- a/src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.h
+++ b/src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhase.h
@@ -1,7 +1,7 @@
 #ifndef B3_GPU_NARROWPHASE_H
 #define B3_GPU_NARROWPHASE_H
 
-#include "Bullet3OpenCL/NarrowphaseCollision/b3Collidable.h"
+#include "Bullet3Collision/NarrowPhaseCollision/shared/b3Collidable.h"
 #include "Bullet3OpenCL/Initialize/b3OpenCLInclude.h"
 #include "Bullet3Common/b3AlignedObjectArray.h"
 #include "Bullet3Common/b3Vector3.h"
@@ -56,7 +56,7 @@ public:
 	void setObjectVelocityCpu(float* linVel, float* angVel, int bodyIndex);
 
 	
-	virtual void computeContacts(cl_mem broadphasePairs, int numBroadphasePairs, cl_mem aabbs, int numObjects);
+	virtual void computeContacts(cl_mem broadphasePairs, int numBroadphasePairs, cl_mem aabbsWorldSpace, int numObjects);
 	
 
 	cl_mem	getBodiesGpu();
@@ -78,7 +78,7 @@ public:
 	cl_mem	getContactsGpu();
 	int	getNumContactsGpu() const;
 
-	cl_mem	getAabbBufferGpu();
+	cl_mem	getAabbLocalSpaceBufferGpu();
 	
 	int getNumRigidBodies() const;
 
diff --git a/src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhaseInternalData.h b/src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhaseInternalData.h
index 20d35388c..14772f20d 100644
--- a/src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhaseInternalData.h
+++ b/src/Bullet3OpenCL/RigidBody/b3GpuNarrowPhaseInternalData.h
@@ -5,7 +5,8 @@
 #include "Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h"
 #include "Bullet3OpenCL/NarrowphaseCollision/b3ConvexPolyhedronCL.h"
 #include "b3Config.h"
-#include "Bullet3OpenCL/NarrowphaseCollision/b3Collidable.h"
+#include "Bullet3Collision/NarrowPhaseCollision/shared/b3Collidable.h"
+
 #include "Bullet3OpenCL/Initialize/b3OpenCLInclude.h"
 #include "Bullet3Common/b3AlignedObjectArray.h"
 #include "Bullet3Common/b3Vector3.h"
@@ -16,8 +17,8 @@
 
 #include "Bullet3OpenCL/NarrowphaseCollision/b3QuantizedBvh.h"
 #include "Bullet3OpenCL/NarrowphaseCollision/b3BvhInfo.h"
-#include "Bullet3Common/b3Int4.h"
-#include "Bullet3Common/b3Int2.h"
+#include "Bullet3Common/shared/b3Int4.h"
+#include "Bullet3Common/shared/b3Int2.h"
 
 
 class b3ConvexUtility;
diff --git a/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.cpp b/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.cpp
index be662b192..cf64ca43a 100644
--- a/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.cpp
+++ b/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.cpp
@@ -450,7 +450,7 @@ void	b3GpuRigidBodyPipeline::setupGpuAabbsFull()
 	launcher.setBuffer(bodies);
 	cl_mem collidables = m_data->m_narrowphase->getCollidablesGpu();
 	launcher.setBuffer(collidables);
-	cl_mem localAabbs = m_data->m_narrowphase->getAabbBufferGpu();
+	cl_mem localAabbs = m_data->m_narrowphase->getAabbLocalSpaceBufferGpu();
 	launcher.setBuffer(localAabbs);
 
 	cl_mem worldAabbs =0;
diff --git a/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipelineInternalData.h b/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipelineInternalData.h
index 312717b93..eafdadb0e 100644
--- a/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipelineInternalData.h
+++ b/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipelineInternalData.h
@@ -20,7 +20,8 @@ subject to the following restrictions:
 #include "Bullet3Common/b3AlignedObjectArray.h"
 
 #include "Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h"
-#include "Bullet3OpenCL/NarrowphaseCollision/b3Collidable.h"
+#include "Bullet3Collision/NarrowPhaseCollision/shared/b3Collidable.h"
+
 
 #include "Bullet3OpenCL/BroadphaseCollision/b3SapAabb.h"
 #include "Bullet3Dynamics/ConstraintSolver/b3TypedConstraint.h"
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.h b/src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.h
index 6fdb72f64..69c9469f4 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.h
+++ b/src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.h
@@ -22,9 +22,13 @@ static const char* batchingKernelsCL= \
 "{\n"
 "	int bla;\n"
 "};\n"
+"#ifdef __cplusplus\n"
+"#else\n"
+"#define b3AtomicInc atomic_inc\n"
+"#endif\n"
 "#endif\n"
 "#ifdef __cplusplus\n"
-"#else//bla\n"
+"#else\n"
 "	typedef float4	b3Float4;\n"
 "#endif \n"
 "#endif //B3_FLOAT4_H\n"
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.h b/src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.h
index b3c1e6e4e..14111699a 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.h
+++ b/src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.h
@@ -22,9 +22,13 @@ static const char* batchingKernelsNewCL= \
 "{\n"
 "	int bla;\n"
 "};\n"
+"#ifdef __cplusplus\n"
+"#else\n"
+"#define b3AtomicInc atomic_inc\n"
+"#endif\n"
 "#endif\n"
 "#ifdef __cplusplus\n"
-"#else//bla\n"
+"#else\n"
 "	typedef float4	b3Float4;\n"
 "#endif \n"
 "#endif //B3_FLOAT4_H\n"
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/solverSetup.h b/src/Bullet3OpenCL/RigidBody/kernels/solverSetup.h
index 05d14c272..d6758e4dc 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/solverSetup.h
+++ b/src/Bullet3OpenCL/RigidBody/kernels/solverSetup.h
@@ -22,9 +22,13 @@ static const char* solverSetupCL= \
 "{\n"
 "	int bla;\n"
 "};\n"
+"#ifdef __cplusplus\n"
+"#else\n"
+"#define b3AtomicInc atomic_inc\n"
+"#endif\n"
 "#endif\n"
 "#ifdef __cplusplus\n"
-"#else//bla\n"
+"#else\n"
 "	typedef float4	b3Float4;\n"
 "#endif \n"
 "#endif //B3_FLOAT4_H\n"
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.h b/src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.h
index 6306685ef..4efa7a1dd 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.h
+++ b/src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.h
@@ -22,9 +22,13 @@ static const char* solverSetup2CL= \
 "{\n"
 "	int bla;\n"
 "};\n"
+"#ifdef __cplusplus\n"
+"#else\n"
+"#define b3AtomicInc atomic_inc\n"
+"#endif\n"
 "#endif\n"
 "#ifdef __cplusplus\n"
-"#else//bla\n"
+"#else\n"
 "	typedef float4	b3Float4;\n"
 "#endif \n"
 "#endif //B3_FLOAT4_H\n"
diff --git a/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h b/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h
index 648c51f69..4f8787bcc 100644
--- a/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h
+++ b/src/Bullet3OpenCL/RigidBody/kernels/solverUtils.h
@@ -22,9 +22,13 @@ static const char* solverUtilsCL= \
 "{\n"
 "	int bla;\n"
 "};\n"
+"#ifdef __cplusplus\n"
+"#else\n"
+"#define b3AtomicInc atomic_inc\n"
+"#endif\n"
 "#endif\n"
 "#ifdef __cplusplus\n"
-"#else//bla\n"
+"#else\n"
 "	typedef float4	b3Float4;\n"
 "#endif \n"
 "#endif //B3_FLOAT4_H\n"
diff --git a/test/OpenCL/BitonicSort/main.cpp b/test/OpenCL/BitonicSort/main.cpp
index dc304209d..643f1ddae 100644
--- a/test/OpenCL/BitonicSort/main.cpp
+++ b/test/OpenCL/BitonicSort/main.cpp
@@ -17,7 +17,7 @@ subject to the following restrictions:
 
 #include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h"
 #include "Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h"
-#include "Bullet3Common/b3Int2.h"
+#include "Bullet3Common/shared/b3Int2.h"
 
 #include "../btgui/Timing/b3Clock.h"
 #include "b3BitonicSort.h"