more rewriting to get this #@!#@!&*( Mac OSX GPU to work

src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.cpp

@@ -14,6 +14,7 @@ subject to the following restrictions:
 */
 
 bool findSeparatingAxisOnGpu = true;
+bool splitSearchSepAxis = false;//true;
 
 bool bvhTraversalKernelGPU = true;
 bool findConcaveSeparatingAxisKernelGPU = true;
@@ -88,6 +89,8 @@ GpuSatCollision::GpuSatCollision(cl_context ctx,cl_device_id device, cl_command_
 m_device(device),
 m_queue(q),
 m_findSeparatingAxisKernel(0),
+m_findSeparatingAxisVertexFaceKernel(0),
+m_findSeparatingAxisEdgeEdgeKernel(0),
 m_totalContactsOut(m_context, m_queue),
 m_sepNormals(m_context, m_queue),
 m_hasSeparatingNormals(m_context, m_queue),
@@ -97,7 +100,8 @@ m_numConcavePairsOut(m_context, m_queue),
 m_gpuCompoundPairs(m_context, m_queue),
 m_gpuCompoundSepNormals(m_context, m_queue),
 m_gpuHasCompoundSepNormals(m_context, m_queue),
-m_numCompoundPairsOut(m_context, m_queue)
+m_numCompoundPairsOut(m_context, m_queue),
+m_dmins(m_context,m_queue)
 {
 	m_totalContactsOut.push_back(0);
 	
@@ -119,6 +123,14 @@ m_numCompoundPairsOut(m_context, m_queue)
 		b3Assert(m_findSeparatingAxisKernel);
 		b3Assert(errNum==CL_SUCCESS);
 
+
+		m_findSeparatingAxisVertexFaceKernel = b3OpenCLUtils::compileCLKernelFromString(m_context, m_device,src, "findSeparatingAxisVertexFaceKernel",&errNum,satProg );
+		b3Assert(m_findSeparatingAxisVertexFaceKernel);
+
+		m_findSeparatingAxisEdgeEdgeKernel = b3OpenCLUtils::compileCLKernelFromString(m_context, m_device,src, "findSeparatingAxisEdgeEdgeKernel",&errNum,satProg );
+		b3Assert(m_findSeparatingAxisVertexFaceKernel);
+
+
 		m_findConcaveSeparatingAxisKernel = b3OpenCLUtils::compileCLKernelFromString(m_context, m_device,src, "findConcaveSeparatingAxisKernel",&errNum,satProg );
 		b3Assert(m_findConcaveSeparatingAxisKernel);
 		b3Assert(errNum==CL_SUCCESS);
@@ -212,6 +224,13 @@ m_numCompoundPairsOut(m_context, m_queue)
 GpuSatCollision::~GpuSatCollision()
 {
 	
+	if (m_findSeparatingAxisVertexFaceKernel)
+		clReleaseKernel(m_findSeparatingAxisVertexFaceKernel);
+
+	if (m_findSeparatingAxisEdgeEdgeKernel)
+		clReleaseKernel(m_findSeparatingAxisEdgeEdgeKernel);
+
+
 	if (m_findSeparatingAxisKernel)
 		clReleaseKernel(m_findSeparatingAxisKernel);
 
@@ -3019,6 +3038,61 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( b3OpenCLArray<b3Int4>*
 		clFinish(m_queue);
 		if (findSeparatingAxisOnGpu)
 		{
+			m_dmins.resize(nPairs);
+			if (splitSearchSepAxis)
+			{
+				{
+				B3_PROFILE("findSeparatingAxisVertexFaceKernel");
+				b3BufferInfoCL bInfo[] = { 
+					b3BufferInfoCL( pairs->getBufferCL(), true ), 
+					b3BufferInfoCL( bodyBuf->getBufferCL(),true), 
+					b3BufferInfoCL( gpuCollidables.getBufferCL(),true), 
+					b3BufferInfoCL( convexData.getBufferCL(),true),
+					b3BufferInfoCL( gpuVertices.getBufferCL(),true),
+					b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
+					b3BufferInfoCL( gpuFaces.getBufferCL(),true),
+					b3BufferInfoCL( gpuIndices.getBufferCL(),true),
+					b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
+					b3BufferInfoCL( m_sepNormals.getBufferCL()),
+					b3BufferInfoCL( m_hasSeparatingNormals.getBufferCL()),
+					b3BufferInfoCL( m_dmins.getBufferCL())
+				};
+
+				b3LauncherCL launcher(m_queue, m_findSeparatingAxisVertexFaceKernel,"findSeparatingAxisVertexFaceKernel");
+				launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
+				launcher.setConst( nPairs  );
+
+				int num = nPairs;
+				launcher.launch1D( num);
+				clFinish(m_queue);
+				}
+				{
+				B3_PROFILE("findSeparatingAxisEdgeEdgeKernel");
+				b3BufferInfoCL bInfo[] = { 
+					b3BufferInfoCL( pairs->getBufferCL(), true ), 
+					b3BufferInfoCL( bodyBuf->getBufferCL(),true), 
+					b3BufferInfoCL( gpuCollidables.getBufferCL(),true), 
+					b3BufferInfoCL( convexData.getBufferCL(),true),
+					b3BufferInfoCL( gpuVertices.getBufferCL(),true),
+					b3BufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
+					b3BufferInfoCL( gpuFaces.getBufferCL(),true),
+					b3BufferInfoCL( gpuIndices.getBufferCL(),true),
+					b3BufferInfoCL( clAabbsWorldSpace.getBufferCL(),true),
+					b3BufferInfoCL( m_sepNormals.getBufferCL()),
+					b3BufferInfoCL( m_hasSeparatingNormals.getBufferCL()),
+					b3BufferInfoCL( m_dmins.getBufferCL())
+				};
+
+				b3LauncherCL launcher(m_queue, m_findSeparatingAxisEdgeEdgeKernel,"findSeparatingAxisEdgeEdgeKernel");
+				launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(b3BufferInfoCL) );
+				launcher.setConst( nPairs  );
+
+				int num = nPairs;
+				launcher.launch1D( num);
+				clFinish(m_queue);
+				}
+
+			} else
 			{
 				B3_PROFILE("findSeparatingAxisKernel");
 				b3BufferInfoCL bInfo[] = { 

src/Bullet3OpenCL/NarrowphaseCollision/b3ConvexHullContact.h

@@ -26,6 +26,10 @@ struct GpuSatCollision
 	cl_device_id			m_device;
 	cl_command_queue		m_queue;
 	cl_kernel				m_findSeparatingAxisKernel;
+
+	cl_kernel m_findSeparatingAxisVertexFaceKernel;
+	cl_kernel m_findSeparatingAxisEdgeEdgeKernel;
+	
 	cl_kernel				m_findConcaveSeparatingAxisKernel;
 	cl_kernel				m_findCompoundPairsKernel;
 	cl_kernel				m_processCompoundPairsKernel;
@@ -50,6 +54,8 @@ struct GpuSatCollision
 	b3OpenCLArray<int>		m_totalContactsOut;
 
 	b3OpenCLArray<b3Vector3> m_sepNormals;
+	b3OpenCLArray<float> m_dmins;
+
 	b3OpenCLArray<int>		m_hasSeparatingNormals;
 	b3OpenCLArray<b3Vector3> m_concaveSepNormals;
 	b3OpenCLArray<int>		m_concaveHasSeparatingNormals;

src/Bullet3OpenCL/NarrowphaseCollision/kernels/sat.cl

@@ -706,7 +706,7 @@ bool findSeparatingAxisEdgeEdge(	__global const ConvexPolyhedronCL* hullA, __glo
 					project(hullB,posB,ornB,&crossje,vertices, &Min1, &Max1);
 				
 					if(Max0<Min1 || Max1<Min0)
-						result = false;
+						return false;
 				
 					float d0 = Max0 - Min1;
 					float d1 = Max1 - Min0;
@@ -1351,6 +1351,176 @@ __kernel void   findSeparatingAxisKernel( __global const int4* pairs,
 }
 
 
+__kernel void   findSeparatingAxisVertexFaceKernel( __global const int4* pairs, 
+																					__global const BodyData* rigidBodies, 
+																					__global const btCollidableGpu* collidables,
+																					__global const ConvexPolyhedronCL* convexShapes, 
+																					__global const float4* vertices,
+																					__global const float4* uniqueEdges,
+																					__global const btGpuFace* faces,
+																					__global const int* indices,
+																					__global btAabbCL* aabbs,
+																					__global volatile float4* separatingNormals,
+																					__global volatile int* hasSeparatingAxis,
+																					__global  float* dmins,
+																					int numPairs
+																					)
+{
+
+	int i = get_global_id(0);
+	
+	if (i<numPairs)
+	{
+
+	
+		int bodyIndexA = pairs[i].x;
+		int bodyIndexB = pairs[i].y;
+
+		int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;
+		int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;
+	
+		int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;
+		int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;
+		
+		
+		//once the broadphase avoids static-static pairs, we can remove this test
+		if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))
+		{
+			hasSeparatingAxis[i] = 0;
+			return;
+		}
+		
+
+		if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONVEX_HULL) ||(collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL))
+		{
+			hasSeparatingAxis[i] = 0;
+			return;
+		}
+			
+		if ((collidables[collidableIndexA].m_shapeType==SHAPE_CONCAVE_TRIMESH))
+		{
+			hasSeparatingAxis[i] = 0;
+			return;
+		}
+
+		int numFacesA = convexShapes[shapeIndexA].m_numFaces;
+
+		float dmin = FLT_MAX;
+
+		dmins[i] = dmin;
+		
+		float4 posA = rigidBodies[bodyIndexA].m_pos;
+		posA.w = 0.f;
+		float4 posB = rigidBodies[bodyIndexB].m_pos;
+		posB.w = 0.f;
+		float4 c0local = convexShapes[shapeIndexA].m_localCenter;
+		float4 ornA = rigidBodies[bodyIndexA].m_quat;
+		float4 c0 = transform(&c0local, &posA, &ornA);
+		float4 c1local = convexShapes[shapeIndexB].m_localCenter;
+		float4 ornB =rigidBodies[bodyIndexB].m_quat;
+		float4 c1 = transform(&c1local,&posB,&ornB);
+		const float4 DeltaC2 = c0 - c1;
+		float4 sepNormal;
+		
+		bool sepA = findSeparatingAxis(	&convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA,
+																								posB,ornB,
+																								DeltaC2,
+																								vertices,uniqueEdges,faces,
+																								indices,&sepNormal,&dmin);
+		hasSeparatingAxis[i] = 4;
+		if (!sepA)
+		{
+			hasSeparatingAxis[i] = 0;
+		} else
+		{
+			bool sepB = findSeparatingAxis(	&convexShapes[shapeIndexB],&convexShapes[shapeIndexA],posB,ornB,
+																									posA,ornA,
+																									DeltaC2,
+																									vertices,uniqueEdges,faces,
+																									indices,&sepNormal,&dmin);
+
+			if (sepB)
+			{
+				dmins[i] = dmin;
+				hasSeparatingAxis[i] = 1;
+				separatingNormals[i] = sepNormal;
+			}
+		}
+		
+	}
+
+}
+
+
+__kernel void   findSeparatingAxisEdgeEdgeKernel( __global const int4* pairs, 
+																					__global const BodyData* rigidBodies, 
+																					__global const btCollidableGpu* collidables,
+																					__global const ConvexPolyhedronCL* convexShapes, 
+																					__global const float4* vertices,
+																					__global const float4* uniqueEdges,
+																					__global const btGpuFace* faces,
+																					__global const int* indices,
+																					__global btAabbCL* aabbs,
+																					__global  float4* separatingNormals,
+																					__global  int* hasSeparatingAxis,
+																					__global  float* dmins,
+																					int numPairs
+																					)
+{
+
+	int i = get_global_id(0);
+	
+	if (i<numPairs)
+	{
+
+		if (hasSeparatingAxis[i])
+		{
+	
+			int bodyIndexA = pairs[i].x;
+			int bodyIndexB = pairs[i].y;
+	
+			int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;
+			int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;
+		
+			int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;
+			int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;
+			
+			
+			int numFacesA = convexShapes[shapeIndexA].m_numFaces;
+	
+			float dmin = dmins[i];
+	
+			float4 posA = rigidBodies[bodyIndexA].m_pos;
+			posA.w = 0.f;
+			float4 posB = rigidBodies[bodyIndexB].m_pos;
+			posB.w = 0.f;
+			float4 c0local = convexShapes[shapeIndexA].m_localCenter;
+			float4 ornA = rigidBodies[bodyIndexA].m_quat;
+			float4 c0 = transform(&c0local, &posA, &ornA);
+			float4 c1local = convexShapes[shapeIndexB].m_localCenter;
+			float4 ornB =rigidBodies[bodyIndexB].m_quat;
+			float4 c1 = transform(&c1local,&posB,&ornB);
+			const float4 DeltaC2 = c0 - c1;
+			float4 sepNormal = separatingNormals[i];
+			
+			bool sepEE = findSeparatingAxisEdgeEdge(	&convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA,
+																									posB,ornB,
+																									DeltaC2,
+																									vertices,uniqueEdges,faces,
+																									indices,&sepNormal,&dmin);
+			if (!sepEE)
+			{
+				hasSeparatingAxis[i] = 0;
+			} else
+			{
+				hasSeparatingAxis[i] = 1;
+				separatingNormals[i] = sepNormal;
+			}
+		}		//if (hasSeparatingAxis[i])
+	}//(i<numPairs)
+}
+
+
 
 
 

src/Bullet3OpenCL/NarrowphaseCollision/kernels/satKernels.h

@@ -901,7 +901,7 @@ static const char* satKernelsCL= \
 "					project(hullB,posB,ornB,&crossje,vertices, &Min1, &Max1);\n"
 "				\n"
 "					if(Max0<Min1 || Max1<Min0)\n"
-"						result = false;\n"
+"						return false;\n"
 "				\n"
 "					float d0 = Max0 - Min1;\n"
 "					float d1 = Max1 - Min0;\n"
@@ -1472,6 +1472,161 @@ static const char* satKernelsCL= \
 "		\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"
+"		\n"
+"		//once the broadphase avoids static-static pairs, we can remove this test\n"
+"		if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0))\n"
+"		{\n"
+"			hasSeparatingAxis[i] = 0;\n"
+"			return;\n"
+"		}\n"
+"		\n"
+"		if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONVEX_HULL) ||(collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL))\n"
+"		{\n"
+"			hasSeparatingAxis[i] = 0;\n"
+"			return;\n"
+"		}\n"
+"			\n"
+"		if ((collidables[collidableIndexA].m_shapeType==SHAPE_CONCAVE_TRIMESH))\n"
+"		{\n"
+"			hasSeparatingAxis[i] = 0;\n"
+"			return;\n"
+"		}\n"
+"		int numFacesA = convexShapes[shapeIndexA].m_numFaces;\n"
+"		float dmin = FLT_MAX;\n"
+"		dmins[i] = dmin;\n"
+"		\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"
+"																					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"
+"			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"
+"		}		//if (hasSeparatingAxis[i])\n"
+"	}//(i<numPairs)\n"
+"}\n"
 "int	findClippingFaces(const float4 separatingNormal,\n"
 "                      const ConvexPolyhedronCL* hullA, \n"
 "					  __global const ConvexPolyhedronCL* hullB,\n"