godot/thirdparty/bullet/Bullet3OpenCL/NarrowphaseCollision/kernels/satKernels.h

//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
static const char* satKernelsCL =
	"//keep this enum in sync with the CPU version (in btCollidable.h)\n"
	"//written by Erwin Coumans\n"
	"#define SHAPE_CONVEX_HULL 3\n"
	"#define SHAPE_CONCAVE_TRIMESH 5\n"
	"#define TRIANGLE_NUM_CONVEX_FACES 5\n"
	"#define SHAPE_COMPOUND_OF_CONVEX_HULLS 6\n"
	"#define B3_MAX_STACK_DEPTH 256\n"
	"typedef unsigned int u32;\n"
	"///keep this in sync with btCollidable.h\n"
	"typedef struct\n"
	"{\n"
	"	union {\n"
	"		int m_numChildShapes;\n"
	"		int m_bvhIndex;\n"
	"	};\n"
	"	union\n"
	"	{\n"
	"		float m_radius;\n"
	"		int	m_compoundBvhIndex;\n"
	"	};\n"
	"	\n"
	"	int m_shapeType;\n"
	"	int m_shapeIndex;\n"
	"	\n"
	"} btCollidableGpu;\n"
	"#define MAX_NUM_PARTS_IN_BITS 10\n"
	"///b3QuantizedBvhNode is a compressed aabb node, 16 bytes.\n"
	"///Node can be used for leafnode or internal node. Leafnodes can point to 32-bit triangle index (non-negative range).\n"
	"typedef struct\n"
	"{\n"
	"	//12 bytes\n"
	"	unsigned short int	m_quantizedAabbMin[3];\n"
	"	unsigned short int	m_quantizedAabbMax[3];\n"
	"	//4 bytes\n"
	"	int	m_escapeIndexOrTriangleIndex;\n"
	"} b3QuantizedBvhNode;\n"
	"typedef struct\n"
	"{\n"
	"	float4		m_aabbMin;\n"
	"	float4		m_aabbMax;\n"
	"	float4		m_quantization;\n"
	"	int			m_numNodes;\n"
	"	int			m_numSubTrees;\n"
	"	int			m_nodeOffset;\n"
	"	int			m_subTreeOffset;\n"
	"} b3BvhInfo;\n"
	"int	getTriangleIndex(const b3QuantizedBvhNode* rootNode)\n"
	"{\n"
	"	unsigned int x=0;\n"
	"	unsigned int y = (~(x&0))<<(31-MAX_NUM_PARTS_IN_BITS);\n"
	"	// Get only the lower bits where the triangle index is stored\n"
	"	return (rootNode->m_escapeIndexOrTriangleIndex&~(y));\n"
	"}\n"
	"int	getTriangleIndexGlobal(__global const b3QuantizedBvhNode* rootNode)\n"
	"{\n"
	"	unsigned int x=0;\n"
	"	unsigned int y = (~(x&0))<<(31-MAX_NUM_PARTS_IN_BITS);\n"
	"	// Get only the lower bits where the triangle index is stored\n"
	"	return (rootNode->m_escapeIndexOrTriangleIndex&~(y));\n"
	"}\n"
	"int isLeafNode(const b3QuantizedBvhNode* rootNode)\n"
	"{\n"
	"	//skipindex is negative (internal node), triangleindex >=0 (leafnode)\n"
	"	return (rootNode->m_escapeIndexOrTriangleIndex >= 0)? 1 : 0;\n"
	"}\n"
	"int isLeafNodeGlobal(__global const b3QuantizedBvhNode* rootNode)\n"
	"{\n"
	"	//skipindex is negative (internal node), triangleindex >=0 (leafnode)\n"
	"	return (rootNode->m_escapeIndexOrTriangleIndex >= 0)? 1 : 0;\n"
	"}\n"
	"	\n"
	"int getEscapeIndex(const b3QuantizedBvhNode* rootNode)\n"
	"{\n"
	"	return -rootNode->m_escapeIndexOrTriangleIndex;\n"
	"}\n"
	"int getEscapeIndexGlobal(__global const b3QuantizedBvhNode* rootNode)\n"
	"{\n"
	"	return -rootNode->m_escapeIndexOrTriangleIndex;\n"
	"}\n"
	"typedef struct\n"
	"{\n"
	"	//12 bytes\n"
	"	unsigned short int	m_quantizedAabbMin[3];\n"
	"	unsigned short int	m_quantizedAabbMax[3];\n"
	"	//4 bytes, points to the root of the subtree\n"
	"	int			m_rootNodeIndex;\n"
	"	//4 bytes\n"
	"	int			m_subtreeSize;\n"
	"	int			m_padding[3];\n"
	"} b3BvhSubtreeInfo;\n"
	"typedef struct\n"
	"{\n"
	"	float4	m_childPosition;\n"
	"	float4	m_childOrientation;\n"
	"	int m_shapeIndex;\n"
	"	int m_unused0;\n"
	"	int m_unused1;\n"
	"	int m_unused2;\n"
	"} btGpuChildShape;\n"
	"typedef struct\n"
	"{\n"
	"	float4 m_pos;\n"
	"	float4 m_quat;\n"
	"	float4 m_linVel;\n"
	"	float4 m_angVel;\n"
	"	u32 m_collidableIdx;\n"
	"	float m_invMass;\n"
	"	float m_restituitionCoeff;\n"
	"	float m_frictionCoeff;\n"
	"} BodyData;\n"
	"typedef struct  \n"
	"{\n"
	"	float4		m_localCenter;\n"
	"	float4		m_extents;\n"
	"	float4		mC;\n"
	"	float4		mE;\n"
	"	\n"
	"	float			m_radius;\n"
	"	int	m_faceOffset;\n"
	"	int m_numFaces;\n"
	"	int	m_numVertices;\n"
	"	int m_vertexOffset;\n"
	"	int	m_uniqueEdgesOffset;\n"
	"	int	m_numUniqueEdges;\n"
	"	int m_unused;\n"
	"} ConvexPolyhedronCL;\n"
	"typedef struct \n"
	"{\n"
	"	union\n"
	"	{\n"
	"		float4	m_min;\n"
	"		float   m_minElems[4];\n"
	"		int			m_minIndices[4];\n"
	"	};\n"
	"	union\n"
	"	{\n"
	"		float4	m_max;\n"
	"		float   m_maxElems[4];\n"
	"		int			m_maxIndices[4];\n"
	"	};\n"
	"} btAabbCL;\n"
	"#ifndef B3_AABB_H\n"
	"#define B3_AABB_H\n"
	"#ifndef B3_FLOAT4_H\n"
	"#define B3_FLOAT4_H\n"
	"#ifndef B3_PLATFORM_DEFINITIONS_H\n"
	"#define B3_PLATFORM_DEFINITIONS_H\n"
	"struct MyTest\n"
	"{\n"
	"	int bla;\n"
	"};\n"
	"#ifdef __cplusplus\n"
	"#else\n"
	"//keep B3_LARGE_FLOAT*B3_LARGE_FLOAT < FLT_MAX\n"
	"#define B3_LARGE_FLOAT 1e18f\n"
	"#define B3_INFINITY 1e18f\n"
	"#define b3Assert(a)\n"
	"#define b3ConstArray(a) __global const a*\n"
	"#define b3AtomicInc atomic_inc\n"
	"#define b3AtomicAdd atomic_add\n"
	"#define b3Fabs fabs\n"
	"#define b3Sqrt native_sqrt\n"
	"#define b3Sin native_sin\n"
	"#define b3Cos native_cos\n"
	"#define B3_STATIC\n"
	"#endif\n"
	"#endif\n"
	"#ifdef __cplusplus\n"
	"#else\n"
	"	typedef float4	b3Float4;\n"
	"	#define b3Float4ConstArg const b3Float4\n"
	"	#define b3MakeFloat4 (float4)\n"
	"	float b3Dot3F4(b3Float4ConstArg v0,b3Float4ConstArg v1)\n"
	"	{\n"
	"		float4 a1 = b3MakeFloat4(v0.xyz,0.f);\n"
	"		float4 b1 = b3MakeFloat4(v1.xyz,0.f);\n"
	"		return dot(a1, b1);\n"
	"	}\n"
	"	b3Float4 b3Cross3(b3Float4ConstArg v0,b3Float4ConstArg v1)\n"
	"	{\n"
	"		float4 a1 = b3MakeFloat4(v0.xyz,0.f);\n"
	"		float4 b1 = b3MakeFloat4(v1.xyz,0.f);\n"
	"		return cross(a1, b1);\n"
	"	}\n"
	"	#define b3MinFloat4 min\n"
	"	#define b3MaxFloat4 max\n"
	"	#define b3Normalized(a) normalize(a)\n"
	"#endif \n"
	"		\n"
	"inline bool b3IsAlmostZero(b3Float4ConstArg v)\n"
	"{\n"
	"	if(b3Fabs(v.x)>1e-6 || b3Fabs(v.y)>1e-6 || b3Fabs(v.z)>1e-6)	\n"
	"		return false;\n"
	"	return true;\n"
	"}\n"
	"inline int    b3MaxDot( b3Float4ConstArg vec, __global const b3Float4* vecArray, int vecLen, float* dotOut )\n"
	"{\n"
	"    float maxDot = -B3_INFINITY;\n"
	"    int i = 0;\n"
	"    int ptIndex = -1;\n"
	"    for( i = 0; i < vecLen; i++ )\n"
	"    {\n"
	"        float dot = b3Dot3F4(vecArray[i],vec);\n"
	"            \n"
	"        if( dot > maxDot )\n"
	"        {\n"
	"            maxDot = dot;\n"
	"            ptIndex = i;\n"
	"        }\n"
	"    }\n"
	"	b3Assert(ptIndex>=0);\n"
	"    if (ptIndex<0)\n"
	"	{\n"
	"		ptIndex = 0;\n"
	"	}\n"
	"    *dotOut = maxDot;\n"
	"    return ptIndex;\n"
	"}\n"
	"#endif //B3_FLOAT4_H\n"
	"#ifndef B3_MAT3x3_H\n"
	"#define B3_MAT3x3_H\n"
	"#ifndef B3_QUAT_H\n"
	"#define B3_QUAT_H\n"
	"#ifndef B3_PLATFORM_DEFINITIONS_H\n"
	"#ifdef __cplusplus\n"
	"#else\n"
	"#endif\n"
	"#endif\n"
	"#ifndef B3_FLOAT4_H\n"
	"#ifdef __cplusplus\n"
	"#else\n"
	"#endif \n"
	"#endif //B3_FLOAT4_H\n"
	"#ifdef __cplusplus\n"
	"#else\n"
	"	typedef float4	b3Quat;\n"
	"	#define b3QuatConstArg const b3Quat\n"
	"	\n"
	"	\n"
	"inline float4 b3FastNormalize4(float4 v)\n"
	"{\n"
	"	v = (float4)(v.xyz,0.f);\n"
	"	return fast_normalize(v);\n"
	"}\n"
	"	\n"
	"inline b3Quat b3QuatMul(b3Quat a, b3Quat b);\n"
	"inline b3Quat b3QuatNormalized(b3QuatConstArg in);\n"
	"inline b3Quat b3QuatRotate(b3QuatConstArg q, b3QuatConstArg vec);\n"
	"inline b3Quat b3QuatInvert(b3QuatConstArg q);\n"
	"inline b3Quat b3QuatInverse(b3QuatConstArg q);\n"
	"inline b3Quat b3QuatMul(b3QuatConstArg a, b3QuatConstArg b)\n"
	"{\n"
	"	b3Quat ans;\n"
	"	ans = b3Cross3( a, b );\n"
	"	ans += a.w*b+b.w*a;\n"
	"//	ans.w = a.w*b.w - (a.x*b.x+a.y*b.y+a.z*b.z);\n"
	"	ans.w = a.w*b.w - b3Dot3F4(a, b);\n"
	"	return ans;\n"
	"}\n"
	"inline b3Quat b3QuatNormalized(b3QuatConstArg in)\n"
	"{\n"
	"	b3Quat q;\n"
	"	q=in;\n"
	"	//return b3FastNormalize4(in);\n"
	"	float len = native_sqrt(dot(q, q));\n"
	"	if(len > 0.f)\n"
	"	{\n"
	"		q *= 1.f / len;\n"
	"	}\n"
	"	else\n"
	"	{\n"
	"		q.x = q.y = q.z = 0.f;\n"
	"		q.w = 1.f;\n"
	"	}\n"
	"	return q;\n"
	"}\n"
	"inline float4 b3QuatRotate(b3QuatConstArg q, b3QuatConstArg vec)\n"
	"{\n"
	"	b3Quat qInv = b3QuatInvert( q );\n"
	"	float4 vcpy = vec;\n"
	"	vcpy.w = 0.f;\n"
	"	float4 out = b3QuatMul(b3QuatMul(q,vcpy),qInv);\n"
	"	return out;\n"
	"}\n"
	"inline b3Quat b3QuatInverse(b3QuatConstArg q)\n"
	"{\n"
	"	return (b3Quat)(-q.xyz, q.w);\n"
	"}\n"
	"inline b3Quat b3QuatInvert(b3QuatConstArg q)\n"
	"{\n"
	"	return (b3Quat)(-q.xyz, q.w);\n"
	"}\n"
	"inline float4 b3QuatInvRotate(b3QuatConstArg q, b3QuatConstArg vec)\n"
	"{\n"
	"	return b3QuatRotate( b3QuatInvert( q ), vec );\n"
	"}\n"
	"inline b3Float4 b3TransformPoint(b3Float4ConstArg point, b3Float4ConstArg translation, b3QuatConstArg  orientation)\n"
	"{\n"
	"	return b3QuatRotate( orientation, point ) + (translation);\n"
	"}\n"
	"	\n"
	"#endif \n"
	"#endif //B3_QUAT_H\n"
	"#ifdef __cplusplus\n"
	"#else\n"
	"typedef struct\n"
	"{\n"
	"	b3Float4 m_row[3];\n"
	"}b3Mat3x3;\n"
	"#define b3Mat3x3ConstArg const b3Mat3x3\n"
	"#define b3GetRow(m,row) (m.m_row[row])\n"
	"inline b3Mat3x3 b3QuatGetRotationMatrix(b3Quat quat)\n"
	"{\n"
	"	b3Float4 quat2 = (b3Float4)(quat.x*quat.x, quat.y*quat.y, quat.z*quat.z, 0.f);\n"
	"	b3Mat3x3 out;\n"
	"	out.m_row[0].x=1-2*quat2.y-2*quat2.z;\n"
	"	out.m_row[0].y=2*quat.x*quat.y-2*quat.w*quat.z;\n"
	"	out.m_row[0].z=2*quat.x*quat.z+2*quat.w*quat.y;\n"
	"	out.m_row[0].w = 0.f;\n"
	"	out.m_row[1].x=2*quat.x*quat.y+2*quat.w*quat.z;\n"
	"	out.m_row[1].y=1-2*quat2.x-2*quat2.z;\n"
	"	out.m_row[1].z=2*quat.y*quat.z-2*quat.w*quat.x;\n"
	"	out.m_row[1].w = 0.f;\n"
	"	out.m_row[2].x=2*quat.x*quat.z-2*quat.w*quat.y;\n"
	"	out.m_row[2].y=2*quat.y*quat.z+2*quat.w*quat.x;\n"
	"	out.m_row[2].z=1-2*quat2.x-2*quat2.y;\n"
	"	out.m_row[2].w = 0.f;\n"
	"	return out;\n"
	"}\n"
	"inline b3Mat3x3 b3AbsoluteMat3x3(b3Mat3x3ConstArg matIn)\n"
	"{\n"
	"	b3Mat3x3 out;\n"
	"	out.m_row[0] = fabs(matIn.m_row[0]);\n"
	"	out.m_row[1] = fabs(matIn.m_row[1]);\n"
	"	out.m_row[2] = fabs(matIn.m_row[2]);\n"
	"	return out;\n"
	"}\n"
	"__inline\n"
	"b3Mat3x3 mtZero();\n"
	"__inline\n"
	"b3Mat3x3 mtIdentity();\n"
	"__inline\n"
	"b3Mat3x3 mtTranspose(b3Mat3x3 m);\n"
	"__inline\n"
	"b3Mat3x3 mtMul(b3Mat3x3 a, b3Mat3x3 b);\n"
	"__inline\n"
	"b3Float4 mtMul1(b3Mat3x3 a, b3Float4 b);\n"
	"__inline\n"
	"b3Float4 mtMul3(b3Float4 a, b3Mat3x3 b);\n"
	"__inline\n"
	"b3Mat3x3 mtZero()\n"
	"{\n"
	"	b3Mat3x3 m;\n"
	"	m.m_row[0] = (b3Float4)(0.f);\n"
	"	m.m_row[1] = (b3Float4)(0.f);\n"
	"	m.m_row[2] = (b3Float4)(0.f);\n"
	"	return m;\n"
	"}\n"
	"__inline\n"
	"b3Mat3x3 mtIdentity()\n"
	"{\n"
	"	b3Mat3x3 m;\n"
	"	m.m_row[0] = (b3Float4)(1,0,0,0);\n"
	"	m.m_row[1] = (b3Float4)(0,1,0,0);\n"
	"	m.m_row[2] = (b3Float4)(0,0,1,0);\n"
	"	return m;\n"
	"}\n"
	"__inline\n"
	"b3Mat3x3 mtTranspose(b3Mat3x3 m)\n"
	"{\n"
	"	b3Mat3x3 out;\n"
	"	out.m_row[0] = (b3Float4)(m.m_row[0].x, m.m_row[1].x, m.m_row[2].x, 0.f);\n"
	"	out.m_row[1] = (b3Float4)(m.m_row[0].y, m.m_row[1].y, m.m_row[2].y, 0.f);\n"
	"	out.m_row[2] = (b3Float4)(m.m_row[0].z, m.m_row[1].z, m.m_row[2].z, 0.f);\n"
	"	return out;\n"
	"}\n"
	"__inline\n"
	"b3Mat3x3 mtMul(b3Mat3x3 a, b3Mat3x3 b)\n"
	"{\n"
	"	b3Mat3x3 transB;\n"
	"	transB = mtTranspose( b );\n"
	"	b3Mat3x3 ans;\n"
	"	//	why this doesn't run when 0ing in the for{}\n"
	"	a.m_row[0].w = 0.f;\n"
	"	a.m_row[1].w = 0.f;\n"
	"	a.m_row[2].w = 0.f;\n"
	"	for(int i=0; i<3; i++)\n"
	"	{\n"
	"//	a.m_row[i].w = 0.f;\n"
	"		ans.m_row[i].x = b3Dot3F4(a.m_row[i],transB.m_row[0]);\n"
	"		ans.m_row[i].y = b3Dot3F4(a.m_row[i],transB.m_row[1]);\n"
	"		ans.m_row[i].z = b3Dot3F4(a.m_row[i],transB.m_row[2]);\n"
	"		ans.m_row[i].w = 0.f;\n"
	"	}\n"
	"	return ans;\n"
	"}\n"
	"__inline\n"
	"b3Float4 mtMul1(b3Mat3x3 a, b3Float4 b)\n"
	"{\n"
	"	b3Float4 ans;\n"
	"	ans.x = b3Dot3F4( a.m_row[0], b );\n"
	"	ans.y = b3Dot3F4( a.m_row[1], b );\n"
	"	ans.z = b3Dot3F4( a.m_row[2], b );\n"
	"	ans.w = 0.f;\n"
	"	return ans;\n"
	"}\n"
	"__inline\n"
	"b3Float4 mtMul3(b3Float4 a, b3Mat3x3 b)\n"
	"{\n"
	"	b3Float4 colx = b3MakeFloat4(b.m_row[0].x, b.m_row[1].x, b.m_row[2].x, 0);\n"
	"	b3Float4 coly = b3MakeFloat4(b.m_row[0].y, b.m_row[1].y, b.m_row[2].y, 0);\n"
	"	b3Float4 colz = b3MakeFloat4(b.m_row[0].z, b.m_row[1].z, b.m_row[2].z, 0);\n"
	"	b3Float4 ans;\n"
	"	ans.x = b3Dot3F4( a, colx );\n"
	"	ans.y = b3Dot3F4( a, coly );\n"
	"	ans.z = b3Dot3F4( a, colz );\n"
	"	return ans;\n"
	"}\n"
	"#endif\n"
	"#endif //B3_MAT3x3_H\n"
	"typedef struct b3Aabb b3Aabb_t;\n"
	"struct b3Aabb\n"
	"{\n"
	"	union\n"
	"	{\n"
	"		float m_min[4];\n"
	"		b3Float4 m_minVec;\n"
	"		int m_minIndices[4];\n"
	"	};\n"
	"	union\n"
	"	{\n"
	"		float	m_max[4];\n"
	"		b3Float4 m_maxVec;\n"
	"		int m_signedMaxIndices[4];\n"
	"	};\n"
	"};\n"
	"inline void b3TransformAabb2(b3Float4ConstArg localAabbMin,b3Float4ConstArg localAabbMax, float margin,\n"
	"						b3Float4ConstArg pos,\n"
	"						b3QuatConstArg orn,\n"
	"						b3Float4* aabbMinOut,b3Float4* aabbMaxOut)\n"
	"{\n"
	"		b3Float4 localHalfExtents = 0.5f*(localAabbMax-localAabbMin);\n"
	"		localHalfExtents+=b3MakeFloat4(margin,margin,margin,0.f);\n"
	"		b3Float4 localCenter = 0.5f*(localAabbMax+localAabbMin);\n"
	"		b3Mat3x3 m;\n"
	"		m = b3QuatGetRotationMatrix(orn);\n"
	"		b3Mat3x3 abs_b = b3AbsoluteMat3x3(m);\n"
	"		b3Float4 center = b3TransformPoint(localCenter,pos,orn);\n"
	"		\n"
	"		b3Float4 extent = b3MakeFloat4(b3Dot3F4(localHalfExtents,b3GetRow(abs_b,0)),\n"
	"										 b3Dot3F4(localHalfExtents,b3GetRow(abs_b,1)),\n"
	"										 b3Dot3F4(localHalfExtents,b3GetRow(abs_b,2)),\n"
	"										 0.f);\n"
	"		*aabbMinOut = center-extent;\n"
	"		*aabbMaxOut = center+extent;\n"
	"}\n"
	"/// conservative test for overlap between two aabbs\n"
	"inline bool b3TestAabbAgainstAabb(b3Float4ConstArg aabbMin1,b3Float4ConstArg aabbMax1,\n"
	"								b3Float4ConstArg aabbMin2, b3Float4ConstArg aabbMax2)\n"
	"{\n"
	"	bool overlap = true;\n"
	"	overlap = (aabbMin1.x > aabbMax2.x || aabbMax1.x < aabbMin2.x) ? false : overlap;\n"
	"	overlap = (aabbMin1.z > aabbMax2.z || aabbMax1.z < aabbMin2.z) ? false : overlap;\n"
	"	overlap = (aabbMin1.y > aabbMax2.y || aabbMax1.y < aabbMin2.y) ? false : overlap;\n"
	"	return overlap;\n"
	"}\n"
	"#endif //B3_AABB_H\n"
	"/*\n"
	"Bullet Continuous Collision Detection and Physics Library\n"
	"Copyright (c) 2003-2013 Erwin Coumans  http://bulletphysics.org\n"
	"This software is provided 'as-is', without any express or implied warranty.\n"
	"In no event will the authors be held liable for any damages arising from the use of this software.\n"
	"Permission is granted to anyone to use this software for any purpose,\n"
	"including commercial applications, and to alter it and redistribute it freely,\n"
	"subject to the following restrictions:\n"
	"1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required.\n"
	"2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.\n"
	"3. This notice may not be removed or altered from any source distribution.\n"
	"*/\n"
	"#ifndef B3_INT2_H\n"
	"#define B3_INT2_H\n"
	"#ifdef __cplusplus\n"
	"#else\n"
	"#define b3UnsignedInt2 uint2\n"
	"#define b3Int2 int2\n"
	"#define b3MakeInt2 (int2)\n"
	"#endif //__cplusplus\n"
	"#endif\n"
	"typedef struct\n"
	"{\n"
	"	float4 m_plane;\n"
	"	int m_indexOffset;\n"
	"	int m_numIndices;\n"
	"} btGpuFace;\n"
	"#define make_float4 (float4)\n"
	"__inline\n"
	"float4 cross3(float4 a, float4 b)\n"
	"{\n"
	"	return cross(a,b);\n"
	"	\n"
	"//	float4 a1 = make_float4(a.xyz,0.f);\n"
	"//	float4 b1 = make_float4(b.xyz,0.f);\n"
	"//	return cross(a1,b1);\n"
	"//float4 c = make_float4(a.y*b.z - a.z*b.y,a.z*b.x - a.x*b.z,a.x*b.y - a.y*b.x,0.f);\n"
	"	\n"
	"	//	float4 c = make_float4(a.y*b.z - a.z*b.y,1.f,a.x*b.y - a.y*b.x,0.f);\n"
	"	\n"
	"	//return c;\n"
	"}\n"
	"__inline\n"
	"float dot3F4(float4 a, float4 b)\n"
	"{\n"
	"	float4 a1 = make_float4(a.xyz,0.f);\n"
	"	float4 b1 = make_float4(b.xyz,0.f);\n"
	"	return dot(a1, b1);\n"
	"}\n"
	"__inline\n"
	"float4 fastNormalize4(float4 v)\n"
	"{\n"
	"	v = make_float4(v.xyz,0.f);\n"
	"	return fast_normalize(v);\n"
	"}\n"
	"///////////////////////////////////////\n"
	"//	Quaternion\n"
	"///////////////////////////////////////\n"
	"typedef float4 Quaternion;\n"
	"__inline\n"
	"Quaternion qtMul(Quaternion a, Quaternion b);\n"
	"__inline\n"
	"Quaternion qtNormalize(Quaternion in);\n"
	"__inline\n"
	"float4 qtRotate(Quaternion q, float4 vec);\n"
	"__inline\n"
	"Quaternion qtInvert(Quaternion q);\n"
	"__inline\n"
	"Quaternion qtMul(Quaternion a, Quaternion b)\n"
	"{\n"
	"	Quaternion ans;\n"
	"	ans = cross3( a, b );\n"
	"	ans += a.w*b+b.w*a;\n"
	"//	ans.w = a.w*b.w - (a.x*b.x+a.y*b.y+a.z*b.z);\n"
	"	ans.w = a.w*b.w - dot3F4(a, b);\n"
	"	return ans;\n"
	"}\n"
	"__inline\n"
	"Quaternion qtNormalize(Quaternion in)\n"
	"{\n"
	"	return fastNormalize4(in);\n"
	"//	in /= length( in );\n"
	"//	return in;\n"
	"}\n"
	"__inline\n"
	"float4 qtRotate(Quaternion q, float4 vec)\n"
	"{\n"
	"	Quaternion qInv = qtInvert( q );\n"
	"	float4 vcpy = vec;\n"
	"	vcpy.w = 0.f;\n"
	"	float4 out = qtMul(qtMul(q,vcpy),qInv);\n"
	"	return out;\n"
	"}\n"
	"__inline\n"
	"Quaternion qtInvert(Quaternion q)\n"
	"{\n"
	"	return (Quaternion)(-q.xyz, q.w);\n"
	"}\n"
	"__inline\n"
	"float4 qtInvRotate(const Quaternion q, float4 vec)\n"
	"{\n"
	"	return qtRotate( qtInvert( q ), vec );\n"
	"}\n"
	"__inline\n"
	"float4 transform(const float4* p, const float4* translation, const Quaternion* orientation)\n"
	"{\n"
	"	return qtRotate( *orientation, *p ) + (*translation);\n"
	"}\n"
	"__inline\n"
	"float4 normalize3(const float4 a)\n"
	"{\n"
	"	float4 n = make_float4(a.x, a.y, a.z, 0.f);\n"
	"	return fastNormalize4( n );\n"
	"}\n"
	"inline void projectLocal(const ConvexPolyhedronCL* hull,  const float4 pos, const float4 orn, \n"
	"const float4* dir, const float4* vertices, float* min, float* max)\n"
	"{\n"
	"	min[0] = FLT_MAX;\n"
	"	max[0] = -FLT_MAX;\n"
	"	int numVerts = hull->m_numVertices;\n"
	"	const float4 localDir = qtInvRotate(orn,*dir);\n"
	"	float offset = dot(pos,*dir);\n"
	"	for(int i=0;i<numVerts;i++)\n"
	"	{\n"
	"		float dp = dot(vertices[hull->m_vertexOffset+i],localDir);\n"
	"		if(dp < min[0])	\n"
	"			min[0] = dp;\n"
	"		if(dp > max[0])	\n"
	"			max[0] = dp;\n"
	"	}\n"
	"	if(min[0]>max[0])\n"
	"	{\n"
	"		float tmp = min[0];\n"
	"		min[0] = max[0];\n"
	"		max[0] = tmp;\n"
	"	}\n"
	"	min[0] += offset;\n"
	"	max[0] += offset;\n"
	"}\n"
	"inline void project(__global const ConvexPolyhedronCL* hull,  const float4 pos, const float4 orn, \n"
	"const float4* dir, __global const float4* vertices, float* min, float* max)\n"
	"{\n"
	"	min[0] = FLT_MAX;\n"
	"	max[0] = -FLT_MAX;\n"
	"	int numVerts = hull->m_numVertices;\n"
	"	const float4 localDir = qtInvRotate(orn,*dir);\n"
	"	float offset = dot(pos,*dir);\n"
	"	for(int i=0;i<numVerts;i++)\n"
	"	{\n"
	"		float dp = dot(vertices[hull->m_vertexOffset+i],localDir);\n"
	"		if(dp < min[0])	\n"
	"			min[0] = dp;\n"
	"		if(dp > max[0])	\n"
	"			max[0] = dp;\n"
	"	}\n"
	"	if(min[0]>max[0])\n"
	"	{\n"
	"		float tmp = min[0];\n"
	"		min[0] = max[0];\n"
	"		max[0] = tmp;\n"
	"	}\n"
	"	min[0] += offset;\n"
	"	max[0] += offset;\n"
	"}\n"
	"inline bool TestSepAxisLocalA(const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n"
	"	const float4 posA,const float4 ornA,\n"
	"	const float4 posB,const float4 ornB,\n"
	"	float4* sep_axis, const float4* verticesA, __global const float4* verticesB,float* depth)\n"
	"{\n"
	"	float Min0,Max0;\n"
	"	float Min1,Max1;\n"
	"	projectLocal(hullA,posA,ornA,sep_axis,verticesA, &Min0, &Max0);\n"
	"	project(hullB,posB,ornB, sep_axis,verticesB, &Min1, &Max1);\n"
	"	if(Max0<Min1 || Max1<Min0)\n"
	"		return false;\n"
	"	float d0 = Max0 - Min1;\n"
	"	float d1 = Max1 - Min0;\n"
	"	*depth = d0<d1 ? d0:d1;\n"
	"	return true;\n"
	"}\n"
	"inline bool IsAlmostZero(const float4 v)\n"
	"{\n"
	"	if(fabs(v.x)>1e-6f || fabs(v.y)>1e-6f || fabs(v.z)>1e-6f)\n"
	"		return false;\n"
	"	return true;\n"
	"}\n"
	"bool findSeparatingAxisLocalA(	const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n"
	"	const float4 posA1,\n"
	"	const float4 ornA,\n"
	"	const float4 posB1,\n"
	"	const float4 ornB,\n"
	"	const float4 DeltaC2,\n"
	"	\n"
	"	const float4* verticesA, \n"
	"	const float4* uniqueEdgesA, \n"
	"	const btGpuFace* facesA,\n"
	"	const int*  indicesA,\n"
	"	__global const float4* verticesB, \n"
	"	__global const float4* uniqueEdgesB, \n"
	"	__global const btGpuFace* facesB,\n"
	"	__global const int*  indicesB,\n"
	"	float4* sep,\n"
	"	float* dmin)\n"
	"{\n"
	"	\n"
	"	float4 posA = posA1;\n"
	"	posA.w = 0.f;\n"
	"	float4 posB = posB1;\n"
	"	posB.w = 0.f;\n"
	"	int curPlaneTests=0;\n"
	"	{\n"
	"		int numFacesA = hullA->m_numFaces;\n"
	"		// Test normals from hullA\n"
	"		for(int i=0;i<numFacesA;i++)\n"
	"		{\n"
	"			const float4 normal = facesA[hullA->m_faceOffset+i].m_plane;\n"
	"			float4 faceANormalWS = qtRotate(ornA,normal);\n"
	"			if (dot3F4(DeltaC2,faceANormalWS)<0)\n"
	"				faceANormalWS*=-1.f;\n"
	"			curPlaneTests++;\n"
	"			float d;\n"
	"			if(!TestSepAxisLocalA( hullA, hullB, posA,ornA,posB,ornB,&faceANormalWS, verticesA, verticesB,&d))\n"
	"				return false;\n"
	"			if(d<*dmin)\n"
	"			{\n"
	"				*dmin = d;\n"
	"				*sep = faceANormalWS;\n"
	"			}\n"
	"		}\n"
	"	}\n"
	"	if((dot3F4(-DeltaC2,*sep))>0.0f)\n"
	"	{\n"
	"		*sep = -(*sep);\n"
	"	}\n"
	"	return true;\n"
	"}\n"
	"bool findSeparatingAxisLocalB(	__global const ConvexPolyhedronCL* hullA,  const ConvexPolyhedronCL* hullB, \n"
	"	const float4 posA1,\n"
	"	const float4 ornA,\n"
	"	const float4 posB1,\n"
	"	const float4 ornB,\n"
	"	const float4 DeltaC2,\n"
	"	__global const float4* verticesA, \n"
	"	__global const float4* uniqueEdgesA, \n"
	"	__global const btGpuFace* facesA,\n"
	"	__global const int*  indicesA,\n"
	"	const float4* verticesB,\n"
	"	const float4* uniqueEdgesB, \n"
	"	const btGpuFace* facesB,\n"
	"	const int*  indicesB,\n"
	"	float4* sep,\n"
	"	float* dmin)\n"
	"{\n"
	"	float4 posA = posA1;\n"
	"	posA.w = 0.f;\n"
	"	float4 posB = posB1;\n"
	"	posB.w = 0.f;\n"
	"	int curPlaneTests=0;\n"
	"	{\n"
	"		int numFacesA = hullA->m_numFaces;\n"
	"		// Test normals from hullA\n"
	"		for(int i=0;i<numFacesA;i++)\n"
	"		{\n"
	"			const float4 normal = facesA[hullA->m_faceOffset+i].m_plane;\n"
	"			float4 faceANormalWS = qtRotate(ornA,normal);\n"
	"			if (dot3F4(DeltaC2,faceANormalWS)<0)\n"
	"				faceANormalWS *= -1.f;\n"
	"			curPlaneTests++;\n"
	"			float d;\n"
	"			if(!TestSepAxisLocalA( hullB, hullA, posB,ornB,posA,ornA, &faceANormalWS, verticesB,verticesA, &d))\n"
	"				return false;\n"
	"			if(d<*dmin)\n"
	"			{\n"
	"				*dmin = d;\n"
	"				*sep = faceANormalWS;\n"
	"			}\n"
	"		}\n"
	"	}\n"
	"	if((dot3F4(-DeltaC2,*sep))>0.0f)\n"
	"	{\n"
	"		*sep = -(*sep);\n"
	"	}\n"
	"	return true;\n"
	"}\n"
	"bool findSeparatingAxisEdgeEdgeLocalA(	const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n"
	"	const float4 posA1,\n"
	"	const float4 ornA,\n"
	"	const float4 posB1,\n"
	"	const float4 ornB,\n"
	"	const float4 DeltaC2,\n"
	"	const float4* verticesA, \n"
	"	const float4* uniqueEdgesA, \n"
	"	const btGpuFace* facesA,\n"
	"	const int*  indicesA,\n"
	"	__global const float4* verticesB, \n"
	"	__global const float4* uniqueEdgesB, \n"
	"	__global const btGpuFace* facesB,\n"
	"	__global const int*  indicesB,\n"
	"		float4* sep,\n"
	"	float* dmin)\n"
	"{\n"
	"	float4 posA = posA1;\n"
	"	posA.w = 0.f;\n"
	"	float4 posB = posB1;\n"
	"	posB.w = 0.f;\n"
	"	int curPlaneTests=0;\n"
	"	int curEdgeEdge = 0;\n"
	"	// Test edges\n"
	"	for(int e0=0;e0<hullA->m_numUniqueEdges;e0++)\n"
	"	{\n"
	"		const float4 edge0 = uniqueEdgesA[hullA->m_uniqueEdgesOffset+e0];\n"
	"		float4 edge0World = qtRotate(ornA,edge0);\n"
	"		for(int e1=0;e1<hullB->m_numUniqueEdges;e1++)\n"
	"		{\n"
	"			const float4 edge1 = uniqueEdgesB[hullB->m_uniqueEdgesOffset+e1];\n"
	"			float4 edge1World = qtRotate(ornB,edge1);\n"
	"			float4 crossje = cross3(edge0World,edge1World);\n"
	"			curEdgeEdge++;\n"
	"			if(!IsAlmostZero(crossje))\n"
	"			{\n"
	"				crossje = normalize3(crossje);\n"
	"				if (dot3F4(DeltaC2,crossje)<0)\n"
	"					crossje *= -1.f;\n"
	"				float dist;\n"
	"				bool result = true;\n"
	"				{\n"
	"					float Min0,Max0;\n"
	"					float Min1,Max1;\n"
	"					projectLocal(hullA,posA,ornA,&crossje,verticesA, &Min0, &Max0);\n"
	"					project(hullB,posB,ornB,&crossje,verticesB, &Min1, &Max1);\n"
	"				\n"
	"					if(Max0<Min1 || Max1<Min0)\n"
	"						result = false;\n"
	"				\n"
	"					float d0 = Max0 - Min1;\n"
	"					float d1 = Max1 - Min0;\n"
	"					dist = d0<d1 ? d0:d1;\n"
	"					result = true;\n"
	"				}\n"
	"				\n"
	"				if(dist<*dmin)\n"
	"				{\n"
	"					*dmin = dist;\n"
	"					*sep = crossje;\n"
	"				}\n"
	"			}\n"
	"		}\n"
	"	}\n"
	"	\n"
	"	if((dot3F4(-DeltaC2,*sep))>0.0f)\n"
	"	{\n"
	"		*sep = -(*sep);\n"
	"	}\n"
	"	return true;\n"
	"}\n"
	"inline bool TestSepAxis(__global const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n"
	"	const float4 posA,const float4 ornA,\n"
	"	const float4 posB,const float4 ornB,\n"
	"	float4* sep_axis, __global const float4* vertices,float* depth)\n"
	"{\n"
	"	float Min0,Max0;\n"
	"	float Min1,Max1;\n"
	"	project(hullA,posA,ornA,sep_axis,vertices, &Min0, &Max0);\n"
	"	project(hullB,posB,ornB, sep_axis,vertices, &Min1, &Max1);\n"
	"	if(Max0<Min1 || Max1<Min0)\n"
	"		return false;\n"
	"	float d0 = Max0 - Min1;\n"
	"	float d1 = Max1 - Min0;\n"
	"	*depth = d0<d1 ? d0:d1;\n"
	"	return true;\n"
	"}\n"
	"bool findSeparatingAxis(	__global const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n"
	"	const float4 posA1,\n"
	"	const float4 ornA,\n"
	"	const float4 posB1,\n"
	"	const float4 ornB,\n"
	"	const float4 DeltaC2,\n"
	"	__global const float4* vertices, \n"
	"	__global const float4* uniqueEdges, \n"
	"	__global const btGpuFace* faces,\n"
	"	__global const int*  indices,\n"
	"	float4* sep,\n"
	"	float* dmin)\n"
	"{\n"
	"	\n"
	"	float4 posA = posA1;\n"
	"	posA.w = 0.f;\n"
	"	float4 posB = posB1;\n"
	"	posB.w = 0.f;\n"
	"	\n"
	"	int curPlaneTests=0;\n"
	"	{\n"
	"		int numFacesA = hullA->m_numFaces;\n"
	"		// Test normals from hullA\n"
	"		for(int i=0;i<numFacesA;i++)\n"
	"		{\n"
	"			const float4 normal = faces[hullA->m_faceOffset+i].m_plane;\n"
	"			float4 faceANormalWS = qtRotate(ornA,normal);\n"
	"	\n"
	"			if (dot3F4(DeltaC2,faceANormalWS)<0)\n"
	"				faceANormalWS*=-1.f;\n"
	"				\n"
	"			curPlaneTests++;\n"
	"	\n"
	"			float d;\n"
	"			if(!TestSepAxis( hullA, hullB, posA,ornA,posB,ornB,&faceANormalWS, vertices,&d))\n"
	"				return false;\n"
	"	\n"
	"			if(d<*dmin)\n"
	"			{\n"
	"				*dmin = d;\n"
	"				*sep = faceANormalWS;\n"
	"			}\n"
	"		}\n"
	"	}\n"
	"		if((dot3F4(-DeltaC2,*sep))>0.0f)\n"
	"		{\n"
	"			*sep = -(*sep);\n"
	"		}\n"
	"	\n"
	"	return true;\n"
	"}\n"
	"bool findSeparatingAxisUnitSphere(	__global const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n"
	"	const float4 posA1,\n"
	"	const float4 ornA,\n"
	"	const float4 posB1,\n"
	"	const float4 ornB,\n"
	"	const float4 DeltaC2,\n"
	"	__global const float4* vertices,\n"
	"	__global const float4* unitSphereDirections,\n"
	"	int numUnitSphereDirections,\n"
	"	float4* sep,\n"
	"	float* dmin)\n"
	"{\n"
	"	\n"
	"	float4 posA = posA1;\n"
	"	posA.w = 0.f;\n"
	"	float4 posB = posB1;\n"
	"	posB.w = 0.f;\n"
	"	int curPlaneTests=0;\n"
	"	int curEdgeEdge = 0;\n"
	"	// Test unit sphere directions\n"
	"	for (int i=0;i<numUnitSphereDirections;i++)\n"
	"	{\n"
	"		float4 crossje;\n"
	"		crossje = unitSphereDirections[i];	\n"
	"		if (dot3F4(DeltaC2,crossje)>0)\n"
	"			crossje *= -1.f;\n"
	"		{\n"
	"			float dist;\n"
	"			bool result = true;\n"
	"			float Min0,Max0;\n"
	"			float Min1,Max1;\n"
	"			project(hullA,posA,ornA,&crossje,vertices, &Min0, &Max0);\n"
	"			project(hullB,posB,ornB,&crossje,vertices, &Min1, &Max1);\n"
	"		\n"
	"			if(Max0<Min1 || Max1<Min0)\n"
	"				return false;\n"
	"		\n"
	"			float d0 = Max0 - Min1;\n"
	"			float d1 = Max1 - Min0;\n"
	"			dist = d0<d1 ? d0:d1;\n"
	"			result = true;\n"
	"	\n"
	"			if(dist<*dmin)\n"
	"			{\n"
	"				*dmin = dist;\n"
	"				*sep = crossje;\n"
	"			}\n"
	"		}\n"
	"	}\n"
	"	\n"
	"	if((dot3F4(-DeltaC2,*sep))>0.0f)\n"
	"	{\n"
	"		*sep = -(*sep);\n"
	"	}\n"
	"	return true;\n"
	"}\n"
	"bool findSeparatingAxisEdgeEdge(	__global const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n"
	"	const float4 posA1,\n"
	"	const float4 ornA,\n"
	"	const float4 posB1,\n"
	"	const float4 ornB,\n"
	"	const float4 DeltaC2,\n"
	"	__global const float4* vertices, \n"
	"	__global const float4* uniqueEdges, \n"
	"	__global const btGpuFace* faces,\n"
	"	__global const int*  indices,\n"
	"	float4* sep,\n"
	"	float* dmin)\n"
	"{\n"
	"	\n"
	"	float4 posA = posA1;\n"
	"	posA.w = 0.f;\n"
	"	float4 posB = posB1;\n"
	"	posB.w = 0.f;\n"
	"	int curPlaneTests=0;\n"
	"	int curEdgeEdge = 0;\n"
	"	// Test edges\n"
	"	for(int e0=0;e0<hullA->m_numUniqueEdges;e0++)\n"
	"	{\n"
	"		const float4 edge0 = uniqueEdges[hullA->m_uniqueEdgesOffset+e0];\n"
	"		float4 edge0World = qtRotate(ornA,edge0);\n"
	"		for(int e1=0;e1<hullB->m_numUniqueEdges;e1++)\n"
	"		{\n"
	"			const float4 edge1 = uniqueEdges[hullB->m_uniqueEdgesOffset+e1];\n"
	"			float4 edge1World = qtRotate(ornB,edge1);\n"
	"			float4 crossje = cross3(edge0World,edge1World);\n"
	"			curEdgeEdge++;\n"
	"			if(!IsAlmostZero(crossje))\n"
	"			{\n"
	"				crossje = normalize3(crossje);\n"
	"				if (dot3F4(DeltaC2,crossje)<0)\n"
	"					crossje*=-1.f;\n"
	"					\n"
	"				float dist;\n"
	"				bool result = true;\n"
	"				{\n"
	"					float Min0,Max0;\n"
	"					float Min1,Max1;\n"
	"					project(hullA,posA,ornA,&crossje,vertices, &Min0, &Max0);\n"
	"					project(hullB,posB,ornB,&crossje,vertices, &Min1, &Max1);\n"
	"				\n"
	"					if(Max0<Min1 || Max1<Min0)\n"
	"						return false;\n"
	"				\n"
	"					float d0 = Max0 - Min1;\n"
	"					float d1 = Max1 - Min0;\n"
	"					dist = d0<d1 ? d0:d1;\n"
	"					result = true;\n"
	"				}\n"
	"				\n"
	"				if(dist<*dmin)\n"
	"				{\n"
	"					*dmin = dist;\n"
	"					*sep = crossje;\n"
	"				}\n"
	"			}\n"
	"		}\n"
	"	}\n"
	"	\n"
	"	if((dot3F4(-DeltaC2,*sep))>0.0f)\n"
	"	{\n"
	"		*sep = -(*sep);\n"
	"	}\n"
	"	return true;\n"
	"}\n"
	"// work-in-progress\n"
	"__kernel void   processCompoundPairsKernel( __global const int4* gpuCompoundPairs,\n"
	"																					__global const BodyData* rigidBodies, \n"
	"																					__global const btCollidableGpu* collidables,\n"
	"																					__global const ConvexPolyhedronCL* convexShapes, \n"
	"																					__global const float4* vertices,\n"
	"																					__global const float4* uniqueEdges,\n"
	"																					__global const btGpuFace* faces,\n"
	"																					__global const int* indices,\n"
	"																					__global btAabbCL* aabbs,\n"
	"																					__global const btGpuChildShape* gpuChildShapes,\n"
	"																					__global volatile float4* gpuCompoundSepNormalsOut,\n"
	"																					__global volatile int* gpuHasCompoundSepNormalsOut,\n"
	"																					int numCompoundPairs\n"
	"																					)\n"
	"{\n"
	"	int i = get_global_id(0);\n"
	"	if (i<numCompoundPairs)\n"
	"	{\n"
	"		int bodyIndexA = gpuCompoundPairs[i].x;\n"
	"		int bodyIndexB = gpuCompoundPairs[i].y;\n"
	"		int childShapeIndexA = gpuCompoundPairs[i].z;\n"
	"		int childShapeIndexB = gpuCompoundPairs[i].w;\n"
	"		\n"
	"		int collidableIndexA = -1;\n"
	"		int collidableIndexB = -1;\n"
	"		\n"
	"		float4 ornA = rigidBodies[bodyIndexA].m_quat;\n"
	"		float4 posA = rigidBodies[bodyIndexA].m_pos;\n"
	"		\n"
	"		float4 ornB = rigidBodies[bodyIndexB].m_quat;\n"
	"		float4 posB = rigidBodies[bodyIndexB].m_pos;\n"
	"							\n"
	"		if (childShapeIndexA >= 0)\n"
	"		{\n"
	"			collidableIndexA = gpuChildShapes[childShapeIndexA].m_shapeIndex;\n"
	"			float4 childPosA = gpuChildShapes[childShapeIndexA].m_childPosition;\n"
	"			float4 childOrnA = gpuChildShapes[childShapeIndexA].m_childOrientation;\n"
	"			float4 newPosA = qtRotate(ornA,childPosA)+posA;\n"
	"			float4 newOrnA = qtMul(ornA,childOrnA);\n"
	"			posA = newPosA;\n"
	"			ornA = newOrnA;\n"
	"		} else\n"
	"		{\n"
	"			collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
	"		}\n"
	"		\n"
	"		if (childShapeIndexB>=0)\n"
	"		{\n"
	"			collidableIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;\n"
	"			float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;\n"
	"			float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;\n"
	"			float4 newPosB = transform(&childPosB,&posB,&ornB);\n"
	"			float4 newOrnB = qtMul(ornB,childOrnB);\n"
	"			posB = newPosB;\n"
	"			ornB = newOrnB;\n"
	"		} else\n"
	"		{\n"
	"			collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;	\n"
	"		}\n"
	"	\n"
	"		gpuHasCompoundSepNormalsOut[i] = 0;\n"
	"	\n"
	"		int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;\n"
	"		int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;\n"
	"	\n"
	"		int shapeTypeA = collidables[collidableIndexA].m_shapeType;\n"
	"		int shapeTypeB = collidables[collidableIndexB].m_shapeType;\n"
	"	\n"
	"		if ((shapeTypeA != SHAPE_CONVEX_HULL) || (shapeTypeB != SHAPE_CONVEX_HULL))\n"
	"		{\n"
	"			return;\n"
	"		}\n"
	"		int hasSeparatingAxis = 5;\n"
	"							\n"
	"		int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
	"		float dmin = FLT_MAX;\n"
	"		posA.w = 0.f;\n"
	"		posB.w = 0.f;\n"
	"		float4 c0local = convexShapes[shapeIndexA].m_localCenter;\n"
	"		float4 c0 = transform(&c0local, &posA, &ornA);\n"
	"		float4 c1local = convexShapes[shapeIndexB].m_localCenter;\n"
	"		float4 c1 = transform(&c1local,&posB,&ornB);\n"
	"		const float4 DeltaC2 = c0 - c1;\n"
	"		float4 sepNormal = make_float4(1,0,0,0);\n"
	"		bool sepA = findSeparatingAxis(	&convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA,posB,ornB,DeltaC2,vertices,uniqueEdges,faces,indices,&sepNormal,&dmin);\n"
	"		hasSeparatingAxis = 4;\n"
	"		if (!sepA)\n"
	"		{\n"
	"			hasSeparatingAxis = 0;\n"
	"		} else\n"
	"		{\n"
	"			bool sepB = findSeparatingAxis(	&convexShapes[shapeIndexB],&convexShapes[shapeIndexA],posB,ornB,posA,ornA,DeltaC2,vertices,uniqueEdges,faces,indices,&sepNormal,&dmin);\n"
	"			if (!sepB)\n"
	"			{\n"
	"				hasSeparatingAxis = 0;\n"
	"			} else//(!sepB)\n"
	"			{\n"
	"				bool sepEE = findSeparatingAxisEdgeEdge(	&convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA,posB,ornB,DeltaC2,vertices,uniqueEdges,faces,indices,&sepNormal,&dmin);\n"
	"				if (sepEE)\n"
	"				{\n"
	"						gpuCompoundSepNormalsOut[i] = sepNormal;//fastNormalize4(sepNormal);\n"
	"						gpuHasCompoundSepNormalsOut[i] = 1;\n"
	"				}//sepEE\n"
	"			}//(!sepB)\n"
	"		}//(!sepA)\n"
	"		\n"
	"		\n"
	"	}\n"
	"		\n"
	"}\n"
	"inline b3Float4 MyUnQuantize(const unsigned short* vecIn, b3Float4 quantization, b3Float4 bvhAabbMin)\n"
	"{\n"
	"		b3Float4 vecOut;\n"
	"		vecOut = b3MakeFloat4(\n"
	"			(float)(vecIn[0]) / (quantization.x),\n"
	"			(float)(vecIn[1]) / (quantization.y),\n"
	"			(float)(vecIn[2]) / (quantization.z),\n"
	"			0.f);\n"
	"		vecOut += bvhAabbMin;\n"
	"		return vecOut;\n"
	"}\n"
	"inline b3Float4 MyUnQuantizeGlobal(__global const unsigned short* vecIn, b3Float4 quantization, b3Float4 bvhAabbMin)\n"
	"{\n"
	"		b3Float4 vecOut;\n"
	"		vecOut = b3MakeFloat4(\n"
	"			(float)(vecIn[0]) / (quantization.x),\n"
	"			(float)(vecIn[1]) / (quantization.y),\n"
	"			(float)(vecIn[2]) / (quantization.z),\n"
	"			0.f);\n"
	"		vecOut += bvhAabbMin;\n"
	"		return vecOut;\n"
	"}\n"
	"// work-in-progress\n"
	"__kernel void   findCompoundPairsKernel( __global const int4* pairs, \n"
	"	__global const BodyData* rigidBodies, \n"
	"	__global const btCollidableGpu* collidables,\n"
	"	__global const ConvexPolyhedronCL* convexShapes, \n"
	"	__global const float4* vertices,\n"
	"	__global const float4* uniqueEdges,\n"
	"	__global const btGpuFace* faces,\n"
	"	__global const int* indices,\n"
	"	__global b3Aabb_t* aabbLocalSpace,\n"
	"	__global const btGpuChildShape* gpuChildShapes,\n"
	"	__global volatile int4* gpuCompoundPairsOut,\n"
	"	__global volatile int* numCompoundPairsOut,\n"
	"	__global const b3BvhSubtreeInfo* subtrees,\n"
	"	__global const b3QuantizedBvhNode* quantizedNodes,\n"
	"	__global const b3BvhInfo* bvhInfos,\n"
	"	int numPairs,\n"
	"	int maxNumCompoundPairsCapacity\n"
	"	)\n"
	"{\n"
	"	int i = get_global_id(0);\n"
	"	if (i<numPairs)\n"
	"	{\n"
	"		int bodyIndexA = pairs[i].x;\n"
	"		int bodyIndexB = pairs[i].y;\n"
	"		int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
	"		int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
	"		int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;\n"
	"		int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;\n"
	"		//once the broadphase avoids static-static pairs, we can remove this test\n"
	"		if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))\n"
	"		{\n"
	"			return;\n"
	"		}\n"
	"		if ((collidables[collidableIndexA].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) &&(collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS))\n"
	"		{\n"
	"			int bvhA = collidables[collidableIndexA].m_compoundBvhIndex;\n"
	"			int bvhB = collidables[collidableIndexB].m_compoundBvhIndex;\n"
	"			int numSubTreesA = bvhInfos[bvhA].m_numSubTrees;\n"
	"			int subTreesOffsetA = bvhInfos[bvhA].m_subTreeOffset;\n"
	"			int subTreesOffsetB = bvhInfos[bvhB].m_subTreeOffset;\n"
	"			int numSubTreesB = bvhInfos[bvhB].m_numSubTrees;\n"
	"			\n"
	"			float4 posA = rigidBodies[bodyIndexA].m_pos;\n"
	"			b3Quat ornA = rigidBodies[bodyIndexA].m_quat;\n"
	"			b3Quat ornB = rigidBodies[bodyIndexB].m_quat;\n"
	"			float4 posB = rigidBodies[bodyIndexB].m_pos;\n"
	"			\n"
	"			for (int p=0;p<numSubTreesA;p++)\n"
	"			{\n"
	"				b3BvhSubtreeInfo subtreeA = subtrees[subTreesOffsetA+p];\n"
	"				//bvhInfos[bvhA].m_quantization\n"
	"				b3Float4 treeAminLocal = MyUnQuantize(subtreeA.m_quantizedAabbMin,bvhInfos[bvhA].m_quantization,bvhInfos[bvhA].m_aabbMin);\n"
	"				b3Float4 treeAmaxLocal = MyUnQuantize(subtreeA.m_quantizedAabbMax,bvhInfos[bvhA].m_quantization,bvhInfos[bvhA].m_aabbMin);\n"
	"				b3Float4 aabbAMinOut,aabbAMaxOut;\n"
	"				float margin=0.f;\n"
	"				b3TransformAabb2(treeAminLocal,treeAmaxLocal, margin,posA,ornA,&aabbAMinOut,&aabbAMaxOut);\n"
	"				\n"
	"				for (int q=0;q<numSubTreesB;q++)\n"
	"				{\n"
	"					b3BvhSubtreeInfo subtreeB = subtrees[subTreesOffsetB+q];\n"
	"					b3Float4 treeBminLocal = MyUnQuantize(subtreeB.m_quantizedAabbMin,bvhInfos[bvhB].m_quantization,bvhInfos[bvhB].m_aabbMin);\n"
	"					b3Float4 treeBmaxLocal = MyUnQuantize(subtreeB.m_quantizedAabbMax,bvhInfos[bvhB].m_quantization,bvhInfos[bvhB].m_aabbMin);\n"
	"					b3Float4 aabbBMinOut,aabbBMaxOut;\n"
	"					float margin=0.f;\n"
	"					b3TransformAabb2(treeBminLocal,treeBmaxLocal, margin,posB,ornB,&aabbBMinOut,&aabbBMaxOut);\n"
	"					\n"
	"					\n"
	"					bool aabbOverlap = b3TestAabbAgainstAabb(aabbAMinOut,aabbAMaxOut,aabbBMinOut,aabbBMaxOut);\n"
	"					if (aabbOverlap)\n"
	"					{\n"
	"						\n"
	"						int startNodeIndexA = subtreeA.m_rootNodeIndex+bvhInfos[bvhA].m_nodeOffset;\n"
	"						int endNodeIndexA = startNodeIndexA+subtreeA.m_subtreeSize;\n"
	"						int startNodeIndexB = subtreeB.m_rootNodeIndex+bvhInfos[bvhB].m_nodeOffset;\n"
	"						int endNodeIndexB = startNodeIndexB+subtreeB.m_subtreeSize;\n"
	"						b3Int2 nodeStack[B3_MAX_STACK_DEPTH];\n"
	"						b3Int2 node0;\n"
	"						node0.x = startNodeIndexA;\n"
	"						node0.y = startNodeIndexB;\n"
	"						int maxStackDepth = B3_MAX_STACK_DEPTH;\n"
	"						int depth=0;\n"
	"						nodeStack[depth++]=node0;\n"
	"						do\n"
	"						{\n"
	"							b3Int2 node = nodeStack[--depth];\n"
	"							b3Float4 aMinLocal = MyUnQuantizeGlobal(quantizedNodes[node.x].m_quantizedAabbMin,bvhInfos[bvhA].m_quantization,bvhInfos[bvhA].m_aabbMin);\n"
	"							b3Float4 aMaxLocal = MyUnQuantizeGlobal(quantizedNodes[node.x].m_quantizedAabbMax,bvhInfos[bvhA].m_quantization,bvhInfos[bvhA].m_aabbMin);\n"
	"							b3Float4 bMinLocal = MyUnQuantizeGlobal(quantizedNodes[node.y].m_quantizedAabbMin,bvhInfos[bvhB].m_quantization,bvhInfos[bvhB].m_aabbMin);\n"
	"							b3Float4 bMaxLocal = MyUnQuantizeGlobal(quantizedNodes[node.y].m_quantizedAabbMax,bvhInfos[bvhB].m_quantization,bvhInfos[bvhB].m_aabbMin);\n"
	"							float margin=0.f;\n"
	"							b3Float4 aabbAMinOut,aabbAMaxOut;\n"
	"							b3TransformAabb2(aMinLocal,aMaxLocal, margin,posA,ornA,&aabbAMinOut,&aabbAMaxOut);\n"
	"							b3Float4 aabbBMinOut,aabbBMaxOut;\n"
	"							b3TransformAabb2(bMinLocal,bMaxLocal, margin,posB,ornB,&aabbBMinOut,&aabbBMaxOut);\n"
	"							\n"
	"							bool nodeOverlap = b3TestAabbAgainstAabb(aabbAMinOut,aabbAMaxOut,aabbBMinOut,aabbBMaxOut);\n"
	"							if (nodeOverlap)\n"
	"							{\n"
	"								bool isLeafA = isLeafNodeGlobal(&quantizedNodes[node.x]);\n"
	"								bool isLeafB = isLeafNodeGlobal(&quantizedNodes[node.y]);\n"
	"								bool isInternalA = !isLeafA;\n"
	"								bool isInternalB = !isLeafB;\n"
	"								//fail, even though it might hit two leaf nodes\n"
	"								if (depth+4>maxStackDepth && !(isLeafA && isLeafB))\n"
	"								{\n"
	"									//printf(\"Error: traversal exceeded maxStackDepth\");\n"
	"									continue;\n"
	"								}\n"
	"								if(isInternalA)\n"
	"								{\n"
	"									int nodeAleftChild = node.x+1;\n"
	"									bool isNodeALeftChildLeaf = isLeafNodeGlobal(&quantizedNodes[node.x+1]);\n"
	"									int nodeArightChild = isNodeALeftChildLeaf? node.x+2 : node.x+1 + getEscapeIndexGlobal(&quantizedNodes[node.x+1]);\n"
	"									if(isInternalB)\n"
	"									{					\n"
	"										int nodeBleftChild = node.y+1;\n"
	"										bool isNodeBLeftChildLeaf = isLeafNodeGlobal(&quantizedNodes[node.y+1]);\n"
	"										int nodeBrightChild = isNodeBLeftChildLeaf? node.y+2 : node.y+1 + getEscapeIndexGlobal(&quantizedNodes[node.y+1]);\n"
	"										nodeStack[depth++] = b3MakeInt2(nodeAleftChild, nodeBleftChild);\n"
	"										nodeStack[depth++] = b3MakeInt2(nodeArightChild, nodeBleftChild);\n"
	"										nodeStack[depth++] = b3MakeInt2(nodeAleftChild, nodeBrightChild);\n"
	"										nodeStack[depth++] = b3MakeInt2(nodeArightChild, nodeBrightChild);\n"
	"									}\n"
	"									else\n"
	"									{\n"
	"										nodeStack[depth++] = b3MakeInt2(nodeAleftChild,node.y);\n"
	"										nodeStack[depth++] = b3MakeInt2(nodeArightChild,node.y);\n"
	"									}\n"
	"								}\n"
	"								else\n"
	"								{\n"
	"									if(isInternalB)\n"
	"									{\n"
	"										int nodeBleftChild = node.y+1;\n"
	"										bool isNodeBLeftChildLeaf = isLeafNodeGlobal(&quantizedNodes[node.y+1]);\n"
	"										int nodeBrightChild = isNodeBLeftChildLeaf? node.y+2 : node.y+1 + getEscapeIndexGlobal(&quantizedNodes[node.y+1]);\n"
	"										nodeStack[depth++] = b3MakeInt2(node.x,nodeBleftChild);\n"
	"										nodeStack[depth++] = b3MakeInt2(node.x,nodeBrightChild);\n"
	"									}\n"
	"									else\n"
	"									{\n"
	"										int compoundPairIdx = atomic_inc(numCompoundPairsOut);\n"
	"										if (compoundPairIdx<maxNumCompoundPairsCapacity)\n"
	"										{\n"
	"											int childShapeIndexA = getTriangleIndexGlobal(&quantizedNodes[node.x]);\n"
	"											int childShapeIndexB = getTriangleIndexGlobal(&quantizedNodes[node.y]);\n"
	"											gpuCompoundPairsOut[compoundPairIdx]  = (int4)(bodyIndexA,bodyIndexB,childShapeIndexA,childShapeIndexB);\n"
	"										}\n"
	"									}\n"
	"								}\n"
	"							}\n"
	"						} while (depth);\n"
	"					}\n"
	"				}\n"
	"			}\n"
	"			\n"
	"			return;\n"
	"		}\n"
	"		if ((collidables[collidableIndexA].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) ||(collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS))\n"
	"		{\n"
	"			if (collidables[collidableIndexA].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) \n"
	"			{\n"
	"				int numChildrenA = collidables[collidableIndexA].m_numChildShapes;\n"
	"				for (int c=0;c<numChildrenA;c++)\n"
	"				{\n"
	"					int childShapeIndexA = collidables[collidableIndexA].m_shapeIndex+c;\n"
	"					int childColIndexA = gpuChildShapes[childShapeIndexA].m_shapeIndex;\n"
	"					float4 posA = rigidBodies[bodyIndexA].m_pos;\n"
	"					float4 ornA = rigidBodies[bodyIndexA].m_quat;\n"
	"					float4 childPosA = gpuChildShapes[childShapeIndexA].m_childPosition;\n"
	"					float4 childOrnA = gpuChildShapes[childShapeIndexA].m_childOrientation;\n"
	"					float4 newPosA = qtRotate(ornA,childPosA)+posA;\n"
	"					float4 newOrnA = qtMul(ornA,childOrnA);\n"
	"					int shapeIndexA = collidables[childColIndexA].m_shapeIndex;\n"
	"					b3Aabb_t aabbAlocal = aabbLocalSpace[shapeIndexA];\n"
	"					float margin = 0.f;\n"
	"					\n"
	"					b3Float4 aabbAMinWS;\n"
	"					b3Float4 aabbAMaxWS;\n"
	"					\n"
	"					b3TransformAabb2(aabbAlocal.m_minVec,aabbAlocal.m_maxVec,margin,\n"
	"						newPosA,\n"
	"						newOrnA,\n"
	"						&aabbAMinWS,&aabbAMaxWS);\n"
	"						\n"
	"					\n"
	"					if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
	"					{\n"
	"						int numChildrenB = collidables[collidableIndexB].m_numChildShapes;\n"
	"						for (int b=0;b<numChildrenB;b++)\n"
	"						{\n"
	"							int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+b;\n"
	"							int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;\n"
	"							float4 ornB = rigidBodies[bodyIndexB].m_quat;\n"
	"							float4 posB = rigidBodies[bodyIndexB].m_pos;\n"
	"							float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;\n"
	"							float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;\n"
	"							float4 newPosB = transform(&childPosB,&posB,&ornB);\n"
	"							float4 newOrnB = qtMul(ornB,childOrnB);\n"
	"							int shapeIndexB = collidables[childColIndexB].m_shapeIndex;\n"
	"							b3Aabb_t aabbBlocal = aabbLocalSpace[shapeIndexB];\n"
	"							\n"
	"							b3Float4 aabbBMinWS;\n"
	"							b3Float4 aabbBMaxWS;\n"
	"							\n"
	"							b3TransformAabb2(aabbBlocal.m_minVec,aabbBlocal.m_maxVec,margin,\n"
	"								newPosB,\n"
	"								newOrnB,\n"
	"								&aabbBMinWS,&aabbBMaxWS);\n"
	"								\n"
	"								\n"
	"							\n"
	"							bool aabbOverlap = b3TestAabbAgainstAabb(aabbAMinWS,aabbAMaxWS,aabbBMinWS,aabbBMaxWS);\n"
	"							if (aabbOverlap)\n"
	"							{\n"
	"								int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
	"								float dmin = FLT_MAX;\n"
	"								float4 posA = newPosA;\n"
	"								posA.w = 0.f;\n"
	"								float4 posB = newPosB;\n"
	"								posB.w = 0.f;\n"
	"								float4 c0local = convexShapes[shapeIndexA].m_localCenter;\n"
	"								float4 ornA = newOrnA;\n"
	"								float4 c0 = transform(&c0local, &posA, &ornA);\n"
	"								float4 c1local = convexShapes[shapeIndexB].m_localCenter;\n"
	"								float4 ornB =newOrnB;\n"
	"								float4 c1 = transform(&c1local,&posB,&ornB);\n"
	"								const float4 DeltaC2 = c0 - c1;\n"
	"								{//\n"
	"									int compoundPairIdx = atomic_inc(numCompoundPairsOut);\n"
	"									if (compoundPairIdx<maxNumCompoundPairsCapacity)\n"
	"									{\n"
	"										gpuCompoundPairsOut[compoundPairIdx]  = (int4)(bodyIndexA,bodyIndexB,childShapeIndexA,childShapeIndexB);\n"
	"									}\n"
	"								}//\n"
	"							}//fi(1)\n"
	"						} //for (int b=0\n"
	"					}//if (collidables[collidableIndexB].\n"
	"					else//if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
	"					{\n"
	"						if (1)\n"
	"						{\n"
	"							int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
	"							float dmin = FLT_MAX;\n"
	"							float4 posA = newPosA;\n"
	"							posA.w = 0.f;\n"
	"							float4 posB = rigidBodies[bodyIndexB].m_pos;\n"
	"							posB.w = 0.f;\n"
	"							float4 c0local = convexShapes[shapeIndexA].m_localCenter;\n"
	"							float4 ornA = newOrnA;\n"
	"							float4 c0 = transform(&c0local, &posA, &ornA);\n"
	"							float4 c1local = convexShapes[shapeIndexB].m_localCenter;\n"
	"							float4 ornB = rigidBodies[bodyIndexB].m_quat;\n"
	"							float4 c1 = transform(&c1local,&posB,&ornB);\n"
	"							const float4 DeltaC2 = c0 - c1;\n"
	"							{\n"
	"								int compoundPairIdx = atomic_inc(numCompoundPairsOut);\n"
	"								if (compoundPairIdx<maxNumCompoundPairsCapacity)\n"
	"								{\n"
	"									gpuCompoundPairsOut[compoundPairIdx] = (int4)(bodyIndexA,bodyIndexB,childShapeIndexA,-1);\n"
	"								}//if (compoundPairIdx<maxNumCompoundPairsCapacity)\n"
	"							}//\n"
	"						}//fi (1)\n"
	"					}//if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
	"				}//for (int b=0;b<numChildrenB;b++)	\n"
	"				return;\n"
	"			}//if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
	"			if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONCAVE_TRIMESH) \n"
	"				&& (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS))\n"
	"			{\n"
	"				int numChildrenB = collidables[collidableIndexB].m_numChildShapes;\n"
	"				for (int b=0;b<numChildrenB;b++)\n"
	"				{\n"
	"					int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+b;\n"
	"					int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;\n"
	"					float4 ornB = rigidBodies[bodyIndexB].m_quat;\n"
	"					float4 posB = rigidBodies[bodyIndexB].m_pos;\n"
	"					float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;\n"
	"					float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;\n"
	"					float4 newPosB = qtRotate(ornB,childPosB)+posB;\n"
	"					float4 newOrnB = qtMul(ornB,childOrnB);\n"
	"					int shapeIndexB = collidables[childColIndexB].m_shapeIndex;\n"
	"					//////////////////////////////////////\n"
	"					if (1)\n"
	"					{\n"
	"						int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
	"						float dmin = FLT_MAX;\n"
	"						float4 posA = rigidBodies[bodyIndexA].m_pos;\n"
	"						posA.w = 0.f;\n"
	"						float4 posB = newPosB;\n"
	"						posB.w = 0.f;\n"
	"						float4 c0local = convexShapes[shapeIndexA].m_localCenter;\n"
	"						float4 ornA = rigidBodies[bodyIndexA].m_quat;\n"
	"						float4 c0 = transform(&c0local, &posA, &ornA);\n"
	"						float4 c1local = convexShapes[shapeIndexB].m_localCenter;\n"
	"						float4 ornB =newOrnB;\n"
	"						float4 c1 = transform(&c1local,&posB,&ornB);\n"
	"						const float4 DeltaC2 = c0 - c1;\n"
	"						{//\n"
	"							int compoundPairIdx = atomic_inc(numCompoundPairsOut);\n"
	"							if (compoundPairIdx<maxNumCompoundPairsCapacity)\n"
	"							{\n"
	"								gpuCompoundPairsOut[compoundPairIdx] = (int4)(bodyIndexA,bodyIndexB,-1,childShapeIndexB);\n"
	"							}//fi (compoundPairIdx<maxNumCompoundPairsCapacity)\n"
	"						}//\n"
	"					}//fi (1)	\n"
	"				}//for (int b=0;b<numChildrenB;b++)\n"
	"				return;\n"
	"			}//if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
	"			return;\n"
	"		}//fi ((collidables[collidableIndexA].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) ||(collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS))\n"
	"	}//i<numPairs\n"
	"}\n"
	"// work-in-progress\n"
	"__kernel void   findSeparatingAxisKernel( __global const int4* pairs, \n"
	"																					__global const BodyData* rigidBodies, \n"
	"																					__global const btCollidableGpu* collidables,\n"
	"																					__global const ConvexPolyhedronCL* convexShapes, \n"
	"																					__global const float4* vertices,\n"
	"																					__global const float4* uniqueEdges,\n"
	"																					__global const btGpuFace* faces,\n"
	"																					__global const int* indices,\n"
	"																					__global btAabbCL* aabbs,\n"
	"																					__global volatile float4* separatingNormals,\n"
	"																					__global volatile int* hasSeparatingAxis,\n"
	"																					int numPairs\n"
	"																					)\n"
	"{\n"
	"	int i = get_global_id(0);\n"
	"	\n"
	"	if (i<numPairs)\n"
	"	{\n"
	"	\n"
	"		int bodyIndexA = pairs[i].x;\n"
	"		int bodyIndexB = pairs[i].y;\n"
	"		int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
	"		int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
	"	\n"
	"		int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;\n"
	"		int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;\n"
	"		\n"
	"		\n"
	"		//once the broadphase avoids static-static pairs, we can remove this test\n"
	"		if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))\n"
	"		{\n"
	"			hasSeparatingAxis[i] = 0;\n"
	"			return;\n"
	"		}\n"
	"		\n"
	"		if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONVEX_HULL) ||(collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL))\n"
	"		{\n"
	"			hasSeparatingAxis[i] = 0;\n"
	"			return;\n"
	"		}\n"
	"			\n"
	"		if ((collidables[collidableIndexA].m_shapeType==SHAPE_CONCAVE_TRIMESH))\n"
	"		{\n"
	"			hasSeparatingAxis[i] = 0;\n"
	"			return;\n"
	"		}\n"
	"		int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
	"		float dmin = FLT_MAX;\n"
	"		float4 posA = rigidBodies[bodyIndexA].m_pos;\n"
	"		posA.w = 0.f;\n"
	"		float4 posB = rigidBodies[bodyIndexB].m_pos;\n"
	"		posB.w = 0.f;\n"
	"		float4 c0local = convexShapes[shapeIndexA].m_localCenter;\n"
	"		float4 ornA = rigidBodies[bodyIndexA].m_quat;\n"
	"		float4 c0 = transform(&c0local, &posA, &ornA);\n"
	"		float4 c1local = convexShapes[shapeIndexB].m_localCenter;\n"
	"		float4 ornB =rigidBodies[bodyIndexB].m_quat;\n"
	"		float4 c1 = transform(&c1local,&posB,&ornB);\n"
	"		const float4 DeltaC2 = c0 - c1;\n"
	"		float4 sepNormal;\n"
	"		\n"
	"		bool sepA = findSeparatingAxis(	&convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA,\n"
	"																								posB,ornB,\n"
	"																								DeltaC2,\n"
	"																								vertices,uniqueEdges,faces,\n"
	"																								indices,&sepNormal,&dmin);\n"
	"		hasSeparatingAxis[i] = 4;\n"
	"		if (!sepA)\n"
	"		{\n"
	"			hasSeparatingAxis[i] = 0;\n"
	"		} else\n"
	"		{\n"
	"			bool sepB = findSeparatingAxis(	&convexShapes[shapeIndexB],&convexShapes[shapeIndexA],posB,ornB,\n"
	"																									posA,ornA,\n"
	"																									DeltaC2,\n"
	"																									vertices,uniqueEdges,faces,\n"
	"																									indices,&sepNormal,&dmin);\n"
	"			if (!sepB)\n"
	"			{\n"
	"				hasSeparatingAxis[i] = 0;\n"
	"			} else\n"
	"			{\n"
	"				bool sepEE = findSeparatingAxisEdgeEdge(	&convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA,\n"
	"																									posB,ornB,\n"
	"																									DeltaC2,\n"
	"																									vertices,uniqueEdges,faces,\n"
	"																									indices,&sepNormal,&dmin);\n"
	"				if (!sepEE)\n"
	"				{\n"
	"					hasSeparatingAxis[i] = 0;\n"
	"				} else\n"
	"				{\n"
	"					hasSeparatingAxis[i] = 1;\n"
	"					separatingNormals[i] = sepNormal;\n"
	"				}\n"
	"			}\n"
	"		}\n"
	"		\n"
	"	}\n"
	"}\n"
	"__kernel void   findSeparatingAxisVertexFaceKernel( __global const int4* pairs, \n"
	"																					__global const BodyData* rigidBodies, \n"
	"																					__global const btCollidableGpu* collidables,\n"
	"																					__global const ConvexPolyhedronCL* convexShapes, \n"
	"																					__global const float4* vertices,\n"
	"																					__global const float4* uniqueEdges,\n"
	"																					__global const btGpuFace* faces,\n"
	"																					__global const int* indices,\n"
	"																					__global btAabbCL* aabbs,\n"
	"																					__global volatile float4* separatingNormals,\n"
	"																					__global volatile int* hasSeparatingAxis,\n"
	"																					__global  float* dmins,\n"
	"																					int numPairs\n"
	"																					)\n"
	"{\n"
	"	int i = get_global_id(0);\n"
	"	\n"
	"	if (i<numPairs)\n"
	"	{\n"
	"	\n"
	"		int bodyIndexA = pairs[i].x;\n"
	"		int bodyIndexB = pairs[i].y;\n"
	"		int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
	"		int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
	"	\n"
	"		int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;\n"
	"		int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;\n"
	"	\n"
	"		hasSeparatingAxis[i] = 0;	\n"
	"		\n"
	"		//once the broadphase avoids static-static pairs, we can remove this test\n"
	"		if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))\n"
	"		{\n"
	"			return;\n"
	"		}\n"
	"		\n"
	"		if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONVEX_HULL) ||(collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL))\n"
	"		{\n"
	"			return;\n"
	"		}\n"
	"			\n"
	"		int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
	"		float dmin = FLT_MAX;\n"
	"		dmins[i] = dmin;\n"
	"		\n"
	"		float4 posA = rigidBodies[bodyIndexA].m_pos;\n"
	"		posA.w = 0.f;\n"
	"		float4 posB = rigidBodies[bodyIndexB].m_pos;\n"
	"		posB.w = 0.f;\n"
	"		float4 c0local = convexShapes[shapeIndexA].m_localCenter;\n"
	"		float4 ornA = rigidBodies[bodyIndexA].m_quat;\n"
	"		float4 c0 = transform(&c0local, &posA, &ornA);\n"
	"		float4 c1local = convexShapes[shapeIndexB].m_localCenter;\n"
	"		float4 ornB =rigidBodies[bodyIndexB].m_quat;\n"
	"		float4 c1 = transform(&c1local,&posB,&ornB);\n"
	"		const float4 DeltaC2 = c0 - c1;\n"
	"		float4 sepNormal;\n"
	"		\n"
	"		bool sepA = findSeparatingAxis(	&convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA,\n"
	"																								posB,ornB,\n"
	"																								DeltaC2,\n"
	"																								vertices,uniqueEdges,faces,\n"
	"																								indices,&sepNormal,&dmin);\n"
	"		hasSeparatingAxis[i] = 4;\n"
	"		if (!sepA)\n"
	"		{\n"
	"			hasSeparatingAxis[i] = 0;\n"
	"		} else\n"
	"		{\n"
	"			bool sepB = findSeparatingAxis(	&convexShapes[shapeIndexB],&convexShapes[shapeIndexA],posB,ornB,\n"
	"																									posA,ornA,\n"
	"																									DeltaC2,\n"
	"																									vertices,uniqueEdges,faces,\n"
	"																									indices,&sepNormal,&dmin);\n"
	"			if (sepB)\n"
	"			{\n"
	"				dmins[i] = dmin;\n"
	"				hasSeparatingAxis[i] = 1;\n"
	"				separatingNormals[i] = sepNormal;\n"
	"			}\n"
	"		}\n"
	"		\n"
	"	}\n"
	"}\n"
	"__kernel void   findSeparatingAxisEdgeEdgeKernel( __global const int4* pairs, \n"
	"																					__global const BodyData* rigidBodies, \n"
	"																					__global const btCollidableGpu* collidables,\n"
	"																					__global const ConvexPolyhedronCL* convexShapes, \n"
	"																					__global const float4* vertices,\n"
	"																					__global const float4* uniqueEdges,\n"
	"																					__global const btGpuFace* faces,\n"
	"																					__global const int* indices,\n"
	"																					__global btAabbCL* aabbs,\n"
	"																					__global  float4* separatingNormals,\n"
	"																					__global  int* hasSeparatingAxis,\n"
	"																					__global  float* dmins,\n"
	"																					__global const float4* unitSphereDirections,\n"
	"																					int numUnitSphereDirections,\n"
	"																					int numPairs\n"
	"																					)\n"
	"{\n"
	"	int i = get_global_id(0);\n"
	"	\n"
	"	if (i<numPairs)\n"
	"	{\n"
	"		if (hasSeparatingAxis[i])\n"
	"		{\n"
	"	\n"
	"			int bodyIndexA = pairs[i].x;\n"
	"			int bodyIndexB = pairs[i].y;\n"
	"	\n"
	"			int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
	"			int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
	"		\n"
	"			int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;\n"
	"			int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;\n"
	"			\n"
	"			\n"
	"			int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
	"	\n"
	"			float dmin = dmins[i];\n"
	"	\n"
	"			float4 posA = rigidBodies[bodyIndexA].m_pos;\n"
	"			posA.w = 0.f;\n"
	"			float4 posB = rigidBodies[bodyIndexB].m_pos;\n"
	"			posB.w = 0.f;\n"
	"			float4 c0local = convexShapes[shapeIndexA].m_localCenter;\n"
	"			float4 ornA = rigidBodies[bodyIndexA].m_quat;\n"
	"			float4 c0 = transform(&c0local, &posA, &ornA);\n"
	"			float4 c1local = convexShapes[shapeIndexB].m_localCenter;\n"
	"			float4 ornB =rigidBodies[bodyIndexB].m_quat;\n"
	"			float4 c1 = transform(&c1local,&posB,&ornB);\n"
	"			const float4 DeltaC2 = c0 - c1;\n"
	"			float4 sepNormal = separatingNormals[i];\n"
	"			\n"
	"			\n"
	"			\n"
	"			bool sepEE = false;\n"
	"			int numEdgeEdgeDirections = convexShapes[shapeIndexA].m_numUniqueEdges*convexShapes[shapeIndexB].m_numUniqueEdges;\n"
	"			if (numEdgeEdgeDirections<=numUnitSphereDirections)\n"
	"			{\n"
	"				sepEE = findSeparatingAxisEdgeEdge(	&convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA,\n"
	"																									posB,ornB,\n"
	"																									DeltaC2,\n"
	"																									vertices,uniqueEdges,faces,\n"
	"																									indices,&sepNormal,&dmin);\n"
	"																									\n"
	"					if (!sepEE)\n"
	"					{\n"
	"						hasSeparatingAxis[i] = 0;\n"
	"					} else\n"
	"					{\n"
	"						hasSeparatingAxis[i] = 1;\n"
	"						separatingNormals[i] = sepNormal;\n"
	"					}\n"
	"			}\n"
	"			/*\n"
	"			///else case is a separate kernel, to make Mac OSX OpenCL compiler happy\n"
	"			else\n"
	"			{\n"
	"				sepEE = findSeparatingAxisUnitSphere(&convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA,\n"
	"																									posB,ornB,\n"
	"																									DeltaC2,\n"
	"																									vertices,unitSphereDirections,numUnitSphereDirections,\n"
	"																									&sepNormal,&dmin);\n"
	"					if (!sepEE)\n"
	"					{\n"
	"						hasSeparatingAxis[i] = 0;\n"
	"					} else\n"
	"					{\n"
	"						hasSeparatingAxis[i] = 1;\n"
	"						separatingNormals[i] = sepNormal;\n"
	"					}\n"
	"			}\n"
	"			*/\n"
	"		}		//if (hasSeparatingAxis[i])\n"
	"	}//(i<numPairs)\n"
	"}\n"
	"inline int	findClippingFaces(const float4 separatingNormal,\n"
	"                      const ConvexPolyhedronCL* hullA, \n"
	"					  __global const ConvexPolyhedronCL* hullB,\n"
	"                      const float4 posA, const Quaternion ornA,const float4 posB, const Quaternion ornB,\n"
	"                       __global float4* worldVertsA1,\n"
	"                      __global float4* worldNormalsA1,\n"
	"                      __global float4* worldVertsB1,\n"
	"                      int capacityWorldVerts,\n"
	"                      const float minDist, float maxDist,\n"
	"					  const float4* verticesA,\n"
	"                      const btGpuFace* facesA,\n"
	"                      const int* indicesA,\n"
	"					  __global const float4* verticesB,\n"
	"                      __global const btGpuFace* facesB,\n"
	"                      __global const int* indicesB,\n"
	"                      __global int4* clippingFaces, int pairIndex)\n"
	"{\n"
	"	int numContactsOut = 0;\n"
	"	int numWorldVertsB1= 0;\n"
	"    \n"
	"    \n"
	"	int closestFaceB=0;\n"
	"	float dmax = -FLT_MAX;\n"
	"    \n"
	"	{\n"
	"		for(int face=0;face<hullB->m_numFaces;face++)\n"
	"		{\n"
	"			const float4 Normal = make_float4(facesB[hullB->m_faceOffset+face].m_plane.x,\n"
	"                                              facesB[hullB->m_faceOffset+face].m_plane.y, facesB[hullB->m_faceOffset+face].m_plane.z,0.f);\n"
	"			const float4 WorldNormal = qtRotate(ornB, Normal);\n"
	"			float d = dot3F4(WorldNormal,separatingNormal);\n"
	"			if (d > dmax)\n"
	"			{\n"
	"				dmax = d;\n"
	"				closestFaceB = face;\n"
	"			}\n"
	"		}\n"
	"	}\n"
	"    \n"
	"	{\n"
	"		const btGpuFace polyB = facesB[hullB->m_faceOffset+closestFaceB];\n"
	"		int numVertices = polyB.m_numIndices;\n"
	"        if (numVertices>capacityWorldVerts)\n"
	"            numVertices = capacityWorldVerts;\n"
	"        \n"
	"		for(int e0=0;e0<numVertices;e0++)\n"
	"		{\n"
	"            if (e0<capacityWorldVerts)\n"
	"            {\n"
	"                const float4 b = verticesB[hullB->m_vertexOffset+indicesB[polyB.m_indexOffset+e0]];\n"
	"                worldVertsB1[pairIndex*capacityWorldVerts+numWorldVertsB1++] = transform(&b,&posB,&ornB);\n"
	"            }\n"
	"		}\n"
	"	}\n"
	"    \n"
	"    int closestFaceA=0;\n"
	"	{\n"
	"		float dmin = FLT_MAX;\n"
	"		for(int face=0;face<hullA->m_numFaces;face++)\n"
	"		{\n"
	"			const float4 Normal = make_float4(\n"
	"                                              facesA[hullA->m_faceOffset+face].m_plane.x,\n"
	"                                              facesA[hullA->m_faceOffset+face].m_plane.y,\n"
	"                                              facesA[hullA->m_faceOffset+face].m_plane.z,\n"
	"                                              0.f);\n"
	"			const float4 faceANormalWS = qtRotate(ornA,Normal);\n"
	"            \n"
	"			float d = dot3F4(faceANormalWS,separatingNormal);\n"
	"			if (d < dmin)\n"
	"			{\n"
	"				dmin = d;\n"
	"				closestFaceA = face;\n"
	"                worldNormalsA1[pairIndex] = faceANormalWS;\n"
	"			}\n"
	"		}\n"
	"	}\n"
	"    \n"
	"    int numVerticesA = facesA[hullA->m_faceOffset+closestFaceA].m_numIndices;\n"
	"    if (numVerticesA>capacityWorldVerts)\n"
	"       numVerticesA = capacityWorldVerts;\n"
	"    \n"
	"	for(int e0=0;e0<numVerticesA;e0++)\n"
	"	{\n"
	"        if (e0<capacityWorldVerts)\n"
	"        {\n"
	"            const float4 a = verticesA[hullA->m_vertexOffset+indicesA[facesA[hullA->m_faceOffset+closestFaceA].m_indexOffset+e0]];\n"
	"            worldVertsA1[pairIndex*capacityWorldVerts+e0] = transform(&a, &posA,&ornA);\n"
	"        }\n"
	"    }\n"
	"    \n"
	"    clippingFaces[pairIndex].x = closestFaceA;\n"
	"    clippingFaces[pairIndex].y = closestFaceB;\n"
	"    clippingFaces[pairIndex].z = numVerticesA;\n"
	"    clippingFaces[pairIndex].w = numWorldVertsB1;\n"
	"    \n"
	"    \n"
	"	return numContactsOut;\n"
	"}\n"
	"// work-in-progress\n"
	"__kernel void   findConcaveSeparatingAxisKernel( __global int4* concavePairs,\n"
	"																					__global const BodyData* rigidBodies,\n"
	"																					__global const btCollidableGpu* collidables,\n"
	"																					__global const ConvexPolyhedronCL* convexShapes, \n"
	"																					__global const float4* vertices,\n"
	"																					__global const float4* uniqueEdges,\n"
	"																					__global const btGpuFace* faces,\n"
	"																					__global const int* indices,\n"
	"																					__global const btGpuChildShape* gpuChildShapes,\n"
	"																					__global btAabbCL* aabbs,\n"
	"																					__global float4* concaveSeparatingNormalsOut,\n"
	"																					__global int* concaveHasSeparatingNormals,\n"
	"																					__global int4* clippingFacesOut,\n"
	"																					__global float4* worldVertsA1GPU,\n"
	"																					__global float4*  worldNormalsAGPU,\n"
	"																					__global float4* worldVertsB1GPU,\n"
	"																					int vertexFaceCapacity,\n"
	"																					int numConcavePairs\n"
	"																					)\n"
	"{\n"
	"	int i = get_global_id(0);\n"
	"	if (i>=numConcavePairs)\n"
	"		return;\n"
	"	concaveHasSeparatingNormals[i] = 0;\n"
	"	int pairIdx = i;\n"
	"	int bodyIndexA = concavePairs[i].x;\n"
	"	int bodyIndexB = concavePairs[i].y;\n"
	"	int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
	"	int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
	"	int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;\n"
	"	int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;\n"
	"	if (collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL&&\n"
	"		collidables[collidableIndexB].m_shapeType!=SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
	"	{\n"
	"		concavePairs[pairIdx].w = -1;\n"
	"		return;\n"
	"	}\n"
	"	int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
	"	int numActualConcaveConvexTests = 0;\n"
	"	\n"
	"	int f = concavePairs[i].z;\n"
	"	\n"
	"	bool overlap = false;\n"
	"	\n"
	"	ConvexPolyhedronCL convexPolyhedronA;\n"
	"	//add 3 vertices of the triangle\n"
	"	convexPolyhedronA.m_numVertices = 3;\n"
	"	convexPolyhedronA.m_vertexOffset = 0;\n"
	"	float4	localCenter = make_float4(0.f,0.f,0.f,0.f);\n"
	"	btGpuFace face = faces[convexShapes[shapeIndexA].m_faceOffset+f];\n"
	"	float4 triMinAabb, triMaxAabb;\n"
	"	btAabbCL triAabb;\n"
	"	triAabb.m_min = make_float4(1e30f,1e30f,1e30f,0.f);\n"
	"	triAabb.m_max = make_float4(-1e30f,-1e30f,-1e30f,0.f);\n"
	"	\n"
	"	float4 verticesA[3];\n"
	"	for (int i=0;i<3;i++)\n"
	"	{\n"
	"		int index = indices[face.m_indexOffset+i];\n"
	"		float4 vert = vertices[convexShapes[shapeIndexA].m_vertexOffset+index];\n"
	"		verticesA[i] = vert;\n"
	"		localCenter += vert;\n"
	"			\n"
	"		triAabb.m_min = min(triAabb.m_min,vert);		\n"
	"		triAabb.m_max = max(triAabb.m_max,vert);		\n"
	"	}\n"
	"	overlap = true;\n"
	"	overlap = (triAabb.m_min.x > aabbs[bodyIndexB].m_max.x || triAabb.m_max.x < aabbs[bodyIndexB].m_min.x) ? false : overlap;\n"
	"	overlap = (triAabb.m_min.z > aabbs[bodyIndexB].m_max.z || triAabb.m_max.z < aabbs[bodyIndexB].m_min.z) ? false : overlap;\n"
	"	overlap = (triAabb.m_min.y > aabbs[bodyIndexB].m_max.y || triAabb.m_max.y < aabbs[bodyIndexB].m_min.y) ? false : overlap;\n"
	"		\n"
	"	if (overlap)\n"
	"	{\n"
	"		float dmin = FLT_MAX;\n"
	"		int hasSeparatingAxis=5;\n"
	"		float4 sepAxis=make_float4(1,2,3,4);\n"
	"		int localCC=0;\n"
	"		numActualConcaveConvexTests++;\n"
	"		//a triangle has 3 unique edges\n"
	"		convexPolyhedronA.m_numUniqueEdges = 3;\n"
	"		convexPolyhedronA.m_uniqueEdgesOffset = 0;\n"
	"		float4 uniqueEdgesA[3];\n"
	"		\n"
	"		uniqueEdgesA[0] = (verticesA[1]-verticesA[0]);\n"
	"		uniqueEdgesA[1] = (verticesA[2]-verticesA[1]);\n"
	"		uniqueEdgesA[2] = (verticesA[0]-verticesA[2]);\n"
	"		convexPolyhedronA.m_faceOffset = 0;\n"
	"                                  \n"
	"		float4 normal = make_float4(face.m_plane.x,face.m_plane.y,face.m_plane.z,0.f);\n"
	"                             \n"
	"		btGpuFace facesA[TRIANGLE_NUM_CONVEX_FACES];\n"
	"		int indicesA[3+3+2+2+2];\n"
	"		int curUsedIndices=0;\n"
	"		int fidx=0;\n"
	"		//front size of triangle\n"
	"		{\n"
	"			facesA[fidx].m_indexOffset=curUsedIndices;\n"
	"			indicesA[0] = 0;\n"
	"			indicesA[1] = 1;\n"
	"			indicesA[2] = 2;\n"
	"			curUsedIndices+=3;\n"
	"			float c = face.m_plane.w;\n"
	"			facesA[fidx].m_plane.x = normal.x;\n"
	"			facesA[fidx].m_plane.y = normal.y;\n"
	"			facesA[fidx].m_plane.z = normal.z;\n"
	"			facesA[fidx].m_plane.w = c;\n"
	"			facesA[fidx].m_numIndices=3;\n"
	"		}\n"
	"		fidx++;\n"
	"		//back size of triangle\n"
	"		{\n"
	"			facesA[fidx].m_indexOffset=curUsedIndices;\n"
	"			indicesA[3]=2;\n"
	"			indicesA[4]=1;\n"
	"			indicesA[5]=0;\n"
	"			curUsedIndices+=3;\n"
	"			float c = dot(normal,verticesA[0]);\n"
	"			float c1 = -face.m_plane.w;\n"
	"			facesA[fidx].m_plane.x = -normal.x;\n"
	"			facesA[fidx].m_plane.y = -normal.y;\n"
	"			facesA[fidx].m_plane.z = -normal.z;\n"
	"			facesA[fidx].m_plane.w = c;\n"
	"			facesA[fidx].m_numIndices=3;\n"
	"		}\n"
	"		fidx++;\n"
	"		bool addEdgePlanes = true;\n"
	"		if (addEdgePlanes)\n"
	"		{\n"
	"			int numVertices=3;\n"
	"			int prevVertex = numVertices-1;\n"
	"			for (int i=0;i<numVertices;i++)\n"
	"			{\n"
	"				float4 v0 = verticesA[i];\n"
	"				float4 v1 = verticesA[prevVertex];\n"
	"                                            \n"
	"				float4 edgeNormal = normalize(cross(normal,v1-v0));\n"
	"				float c = -dot(edgeNormal,v0);\n"
	"				facesA[fidx].m_numIndices = 2;\n"
	"				facesA[fidx].m_indexOffset=curUsedIndices;\n"
	"				indicesA[curUsedIndices++]=i;\n"
	"				indicesA[curUsedIndices++]=prevVertex;\n"
	"                                            \n"
	"				facesA[fidx].m_plane.x = edgeNormal.x;\n"
	"				facesA[fidx].m_plane.y = edgeNormal.y;\n"
	"				facesA[fidx].m_plane.z = edgeNormal.z;\n"
	"				facesA[fidx].m_plane.w = c;\n"
	"				fidx++;\n"
	"				prevVertex = i;\n"
	"			}\n"
	"		}\n"
	"		convexPolyhedronA.m_numFaces = TRIANGLE_NUM_CONVEX_FACES;\n"
	"		convexPolyhedronA.m_localCenter = localCenter*(1.f/3.f);\n"
	"		float4 posA = rigidBodies[bodyIndexA].m_pos;\n"
	"		posA.w = 0.f;\n"
	"		float4 posB = rigidBodies[bodyIndexB].m_pos;\n"
	"		posB.w = 0.f;\n"
	"		float4 ornA = rigidBodies[bodyIndexA].m_quat;\n"
	"		float4 ornB =rigidBodies[bodyIndexB].m_quat;\n"
	"		\n"
	"		///////////////////\n"
	"		///compound shape support\n"
	"		if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)\n"
	"		{\n"
	"			int compoundChild = concavePairs[pairIdx].w;\n"
	"			int childShapeIndexB = compoundChild;//collidables[collidableIndexB].m_shapeIndex+compoundChild;\n"
	"			int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex;\n"
	"			float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition;\n"
	"			float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation;\n"
	"			float4 newPosB = transform(&childPosB,&posB,&ornB);\n"
	"			float4 newOrnB = qtMul(ornB,childOrnB);\n"
	"			posB = newPosB;\n"
	"			ornB = newOrnB;\n"
	"			shapeIndexB = collidables[childColIndexB].m_shapeIndex;\n"
	"		}\n"
	"		//////////////////\n"
	"		float4 c0local = convexPolyhedronA.m_localCenter;\n"
	"		float4 c0 = transform(&c0local, &posA, &ornA);\n"
	"		float4 c1local = convexShapes[shapeIndexB].m_localCenter;\n"
	"		float4 c1 = transform(&c1local,&posB,&ornB);\n"
	"		const float4 DeltaC2 = c0 - c1;\n"
	"		bool sepA = findSeparatingAxisLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],\n"
	"												posA,ornA,\n"
	"												posB,ornB,\n"
	"												DeltaC2,\n"
	"												verticesA,uniqueEdgesA,facesA,indicesA,\n"
	"												vertices,uniqueEdges,faces,indices,\n"
	"												&sepAxis,&dmin);\n"
	"		hasSeparatingAxis = 4;\n"
	"		if (!sepA)\n"
	"		{\n"
	"			hasSeparatingAxis = 0;\n"
	"		} else\n"
	"		{\n"
	"			bool sepB = findSeparatingAxisLocalB(	&convexShapes[shapeIndexB],&convexPolyhedronA,\n"
	"												posB,ornB,\n"
	"												posA,ornA,\n"
	"												DeltaC2,\n"
	"												vertices,uniqueEdges,faces,indices,\n"
	"												verticesA,uniqueEdgesA,facesA,indicesA,\n"
	"												&sepAxis,&dmin);\n"
	"			if (!sepB)\n"
	"			{\n"
	"				hasSeparatingAxis = 0;\n"
	"			} else\n"
	"			{\n"
	"				bool sepEE = findSeparatingAxisEdgeEdgeLocalA(	&convexPolyhedronA, &convexShapes[shapeIndexB],\n"
	"															posA,ornA,\n"
	"															posB,ornB,\n"
	"															DeltaC2,\n"
	"															verticesA,uniqueEdgesA,facesA,indicesA,\n"
	"															vertices,uniqueEdges,faces,indices,\n"
	"															&sepAxis,&dmin);\n"
	"	\n"
	"				if (!sepEE)\n"
	"				{\n"
	"					hasSeparatingAxis = 0;\n"
	"				} else\n"
	"				{\n"
	"					hasSeparatingAxis = 1;\n"
	"				}\n"
	"			}\n"
	"		}	\n"
	"		\n"
	"		if (hasSeparatingAxis)\n"
	"		{\n"
	"			sepAxis.w = dmin;\n"
	"			concaveSeparatingNormalsOut[pairIdx]=sepAxis;\n"
	"			concaveHasSeparatingNormals[i]=1;\n"
	"			float minDist = -1e30f;\n"
	"			float maxDist = 0.02f;\n"
	"		\n"
	"			findClippingFaces(sepAxis,\n"
	"                     &convexPolyhedronA,\n"
	"					 &convexShapes[shapeIndexB],\n"
	"					 posA,ornA,\n"
	"					 posB,ornB,\n"
	"                      worldVertsA1GPU,\n"
	"                      worldNormalsAGPU,\n"
	"                      worldVertsB1GPU,\n"
	"					  vertexFaceCapacity,\n"
	"                      minDist, maxDist,\n"
	"                      verticesA,\n"
	"                      facesA,\n"
	"                      indicesA,\n"
	" 					  vertices,\n"
	"                      faces,\n"
	"                      indices,\n"
	"                      clippingFacesOut, pairIdx);\n"
	"		} else\n"
	"		{	\n"
	"			//mark this pair as in-active\n"
	"			concavePairs[pairIdx].w = -1;\n"
	"		}\n"
	"	}\n"
	"	else\n"
	"	{	\n"
	"		//mark this pair as in-active\n"
	"		concavePairs[pairIdx].w = -1;\n"
	"	}\n"
	"	\n"
	"	concavePairs[pairIdx].z = -1;//now z is used for existing/persistent contacts\n"
	"}\n";