From ec0d2ed52311fab1059d8499675fb211249c15cc Mon Sep 17 00:00:00 2001
From: erwin coumans <erwin.coumans@gmail.com>
Date: Fri, 12 Apr 2013 10:25:11 -0700
Subject: [PATCH] extend Wavefront obj loader to recognize objects (for
 compound shape creation) add reproduction of a bug in compound versus concave
 trimesh

---
 data/cube.mtl                                 |  11 +
 data/cube.obj                                 |  26 ++
 demo/ObjLoader/objLoader.cpp                  |   3 +
 demo/ObjLoader/objLoader.h                    |   2 +
 demo/ObjLoader/obj_parser.cpp                 |  29 +-
 demo/ObjLoader/obj_parser.h                   |  16 ++
 demo/gpudemo/GpuDemo.h                        |   6 +-
 demo/gpudemo/main_opengl3core.cpp             |  11 +-
 demo/gpudemo/rigidbody/ConcaveScene.cpp       | 249 +++++++++++++++++-
 demo/gpudemo/rigidbody/ConcaveScene.h         |  21 ++
 opencl/gpu_sat/host/ConvexHullContact.cpp     |  25 +-
 opencl/gpu_sat/kernels/satClipHullContacts.cl |   4 +-
 opencl/gpu_sat/kernels/satClipHullContacts.h  |   4 +-
 13 files changed, 386 insertions(+), 21 deletions(-)
 create mode 100644 data/cube.mtl
 create mode 100644 data/cube.obj

diff --git a/data/cube.mtl b/data/cube.mtl
new file mode 100644
index 000000000..8ce4b8207
--- /dev/null
+++ b/data/cube.mtl
@@ -0,0 +1,11 @@
+# Blender MTL File: 'None'
+# Material Count: 1
+
+newmtl Material
+Ns 96.078431
+Ka 0.000000 0.000000 0.000000
+Kd 0.640000 0.640000 0.640000
+Ks 0.500000 0.500000 0.500000
+Ni 1.000000
+d 1.000000
+illum 2
diff --git a/data/cube.obj b/data/cube.obj
new file mode 100644
index 000000000..5270c7374
--- /dev/null
+++ b/data/cube.obj
@@ -0,0 +1,26 @@
+# Blender v2.66 (sub 1) OBJ File: ''
+# www.blender.org
+mtllib cube.mtl
+o Cube
+v 1.000000 -1.000000 -1.000000
+v 1.000000 -1.000000 1.000000
+v -1.000000 -1.000000 1.000000
+v -1.000000 -1.000000 -1.000000
+v 1.000000 1.000000 -0.999999
+v 0.999999 1.000000 1.000001
+v -1.000000 1.000000 1.000000
+v -1.000000 1.000000 -1.000000
+usemtl Material
+s off
+f 1 2 3
+f 5 8 7
+f 1 5 6
+f 2 6 3
+f 3 7 4
+f 5 1 4
+f 4 1 3
+f 6 5 7
+f 2 1 6
+f 6 7 3
+f 7 8 4
+f 8 5 4
diff --git a/demo/ObjLoader/objLoader.cpp b/demo/ObjLoader/objLoader.cpp
index e7ed84e9c..6816fc782 100644
--- a/demo/ObjLoader/objLoader.cpp
+++ b/demo/ObjLoader/objLoader.cpp
@@ -34,6 +34,9 @@ int objLoader::load(char *filename)
 		this->lightDiscList = data.light_disc_list;
 		this->lightQuadList = data.light_quad_list;
 
+		this->objectList = data.object_list;
+		this->objectCount = data.object_count;
+
 		this->materialList = data.material_list;
 
 		this->camera = data.camera;
diff --git a/demo/ObjLoader/objLoader.h b/demo/ObjLoader/objLoader.h
index 1faa312a8..fd7f976b3 100644
--- a/demo/ObjLoader/objLoader.h
+++ b/demo/ObjLoader/objLoader.h
@@ -26,6 +26,7 @@ public:
 	obj_light_quad **lightQuadList;
 	obj_light_disc **lightDiscList;
 	
+	obj_object** objectList;
 	obj_material **materialList;
 	
 	int vertexCount;
@@ -40,6 +41,7 @@ public:
 	int lightQuadCount;
 	int lightDiscCount;
 
+	int objectCount;
 	int materialCount;
 
 	obj_camera *camera;
diff --git a/demo/ObjLoader/obj_parser.cpp b/demo/ObjLoader/obj_parser.cpp
index e39ddd41d..85fa01b6c 100644
--- a/demo/ObjLoader/obj_parser.cpp
+++ b/demo/ObjLoader/obj_parser.cpp
@@ -89,6 +89,18 @@ int obj_parse_vertex_index(int *vertex_index, int *texture_index, int *normal_in
 	return vertex_count;
 }
 
+
+obj_object* obj_parse_object(obj_growable_scene_data *scene)
+{
+	obj_object* obj= (obj_object*)malloc(sizeof(obj_object));
+	obj->vertex_offset = scene->vertex_list.item_count;
+	obj->face_offset = scene->face_list.item_count;
+	// get the name
+	strncpy(obj->name, strtok(NULL, " \t"), OBJECT_NAME_SIZE);
+	return obj;
+}
+
+
 obj_face* obj_parse_face(obj_growable_scene_data *scene)
 {
 	int vertex_count;
@@ -417,7 +429,11 @@ int obj_parse_obj_file(obj_growable_scene_data *growable_data, char *filename)
 		}
 		
 		else if( strequal(current_token, "o") ) //object name
-		{ }
+		{ 
+			obj_object* obj = obj_parse_object(growable_data);
+			list_add_item(&growable_data->object_list, obj, NULL);
+
+		}
 		else if( strequal(current_token, "s") ) //smoothing
 		{ }
 		else if( strequal(current_token, "g") ) // group
@@ -450,6 +466,8 @@ void obj_init_temp_storage(obj_growable_scene_data *growable_data)
 	list_make(&growable_data->light_quad_list, 10, 1);
 	list_make(&growable_data->light_disc_list, 10, 1);
 	
+	list_make(&growable_data->object_list,10,1);
+
 	list_make(&growable_data->material_list, 10, 1);	
 	
 	growable_data->camera = NULL;
@@ -469,6 +487,7 @@ void obj_free_temp_storage(obj_growable_scene_data *growable_data)
 	obj_free_half_list(&growable_data->light_quad_list);
 	obj_free_half_list(&growable_data->light_disc_list);
 	
+	obj_free_half_list(&growable_data->object_list);
 	obj_free_half_list(&growable_data->material_list);
 }
 
@@ -506,6 +525,11 @@ void delete_obj_data(obj_scene_data *data_out)
 		free(data_out->light_quad_list[i]);
 	free(data_out->light_quad_list);
 
+	for(i=0; i<data_out->object_count; i++)
+		free(data_out->object_list[i]);
+	free(data_out->object_list);
+
+
 	for(i=0; i<data_out->material_count; i++)
 		free(data_out->material_list[i]);
 	free(data_out->material_list);
@@ -541,6 +565,9 @@ void obj_copy_to_out_storage(obj_scene_data *data_out, obj_growable_scene_data *
 	data_out->light_disc_list = (obj_light_disc**)growable_data->light_disc_list.items;
 	data_out->light_quad_list = (obj_light_quad**)growable_data->light_quad_list.items;
 	
+	data_out->object_list = (obj_object**)growable_data->object_list.items;
+	data_out->object_count = growable_data->object_list.item_count;
+
 	data_out->material_list = (obj_material**)growable_data->material_list.items;
 	
 	data_out->camera = growable_data->camera;
diff --git a/demo/ObjLoader/obj_parser.h b/demo/ObjLoader/obj_parser.h
index b9fe34cf5..0ed5d2a8d 100644
--- a/demo/ObjLoader/obj_parser.h
+++ b/demo/ObjLoader/obj_parser.h
@@ -5,9 +5,13 @@
 
 #define OBJ_FILENAME_LENGTH 500
 #define MATERIAL_NAME_SIZE 255
+#define OBJECT_NAME_SIZE 255
+
 #define OBJ_LINE_SIZE 500
 #define MAX_VERTEX_COUNT 4 //can only handle quads or triangles
 
+
+
 typedef struct obj_face
 {
 	int vertex_index[MAX_VERTEX_COUNT];
@@ -81,6 +85,13 @@ typedef struct obj_light_quad
 	int material_index;
 };
 
+typedef struct obj_object
+{
+	int vertex_offset;
+	int face_offset;
+	char name[OBJECT_NAME_SIZE];
+};
+
 typedef struct obj_growable_scene_data
 {
 //	vector extreme_dimensions[2];
@@ -100,6 +111,8 @@ typedef struct obj_growable_scene_data
 	list light_disc_list;
 	
 	list material_list;
+
+	list object_list;
 	
 	obj_camera *camera;
 };
@@ -120,6 +133,9 @@ typedef struct obj_scene_data
 	
 	obj_material **material_list;
 	
+	obj_object** object_list;
+	int object_count;
+
 	int vertex_count;
 	int vertex_normal_count;
 	int vertex_texture_count;
diff --git a/demo/gpudemo/GpuDemo.h b/demo/gpudemo/GpuDemo.h
index 4b34bf868..c8ef1303c 100644
--- a/demo/gpudemo/GpuDemo.h
+++ b/demo/gpudemo/GpuDemo.h
@@ -38,9 +38,9 @@ public:
                     :useOpenCL(true),
                     preferredOpenCLPlatformIndex(-1),
                     preferredOpenCLDeviceIndex(-1),
-					arraySizeX(10),
-		arraySizeY(30),
-		arraySizeZ(10),
+					arraySizeX(1),
+		arraySizeY(2),
+		arraySizeZ(1),
 		m_useConcaveMesh(false),
 		gapX(14.3),
 		gapY(14.0),
diff --git a/demo/gpudemo/main_opengl3core.cpp b/demo/gpudemo/main_opengl3core.cpp
index 324b85e5a..313fe052c 100644
--- a/demo/gpudemo/main_opengl3core.cpp
+++ b/demo/gpudemo/main_opengl3core.cpp
@@ -65,15 +65,20 @@ btAlignedObjectArray<const char*> demoNames;
 int selectedDemo = 0;
 GpuDemo::CreateFunc* allDemos[]=
 {
+		ConcaveCompound2Scene::MyCreateFunc,
+
+	ConcaveCompoundScene::MyCreateFunc,
+
+	GpuCompoundPlaneScene::MyCreateFunc,
+
+
 	GpuBoxPlaneScene::MyCreateFunc,	
 	GpuConvexPlaneScene::MyCreateFunc,
 	
 
 	GpuCompoundScene::MyCreateFunc,
-		GpuCompoundPlaneScene::MyCreateFunc,
-
+	
 	GpuConvexScene::MyCreateFunc,
-	ConcaveCompoundScene::MyCreateFunc,
 
 	ConcaveSphereScene::MyCreateFunc,
 
diff --git a/demo/gpudemo/rigidbody/ConcaveScene.cpp b/demo/gpudemo/rigidbody/ConcaveScene.cpp
index 58948575d..34ebc95ec 100644
--- a/demo/gpudemo/rigidbody/ConcaveScene.cpp
+++ b/demo/gpudemo/rigidbody/ConcaveScene.cpp
@@ -160,11 +160,13 @@ void ConcaveScene::createConcaveMesh(const ConstructionInfo& ci)
 	objLoader* objData = new objLoader();
 	//char* fileName = "data/slopedPlane100.obj";
 	//char* fileName = "data/plane100.obj";
+	char* fileName = "data/plane100.obj";
 
 	//char* fileName = "data/teddy.obj";//"plane.obj";
 //	char* fileName = "data/sponza_closed.obj";//"plane.obj";
 	//char* fileName = "data/leoTest1.obj";
-	char* fileName = "data/samurai_monastry.obj";
+	//char* fileName = "data/samurai_monastry.obj";
+//	char* fileName = "data/teddy2_VHACD_CHs.obj";
 	
 	btVector3 shift(0,0,0);//0,230,80);//150,-100,-120);
 	btVector4 scaling(4,4,4,1);
@@ -242,13 +244,39 @@ void ConcaveScene::createConcaveMesh(const ConstructionInfo& ci)
 
 		}
 	}
+	
+	delete objData;
 
 }
 
 void ConcaveScene::setupScene(const ConstructionInfo& ci)
 {
 
-	createConcaveMesh(ci);
+	if (1)
+	{
+		createConcaveMesh(ci);
+	} else
+	{
+		int strideInBytes = 9*sizeof(float);
+		int numVertices = sizeof(cube_vertices)/strideInBytes;
+		int numIndices = sizeof(cube_indices)/sizeof(int);
+		int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
+		int group=1;
+		int mask=1;
+		int index=0;
+		{
+			btVector4 scaling(400,0.001,400,1);
+			int colIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
+			btVector3 position(0,-2,0);
+			btQuaternion orn(0,0,0,1);
+				
+			btVector4 color(0,0,1,1);
+		
+			int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
+			int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(0.f,position,orn,colIndex,index);
+	
+		}
+	}
 
 	createDynamicObjects(ci);
 	
@@ -256,7 +284,7 @@ void ConcaveScene::setupScene(const ConstructionInfo& ci)
 	//float camPos[4]={1,12.5,1.5,0};
 	m_instancingRenderer->setCameraPitch(45);
 	m_instancingRenderer->setCameraTargetPosition(camPos);
-	m_instancingRenderer->setCameraDistance(370);
+	m_instancingRenderer->setCameraDistance(25);
 
 }
 
@@ -327,6 +355,221 @@ void ConcaveCompoundScene::setupScene(const ConstructionInfo& ci)
 
 }
 
+void ConcaveCompound2Scene::createDynamicObjects(const ConstructionInfo& ci)
+{
+	objLoader* objData = new objLoader();
+
+	//char* fileName = "data/teddy2_VHACD_CHs.obj";
+	char* fileName = "data/cube_offset.obj";
+
+	
+	btVector3 shift(0,0,0);//0,230,80);//150,-100,-120);
+	btVector4 scaling(1,1,1,1);
+	FILE* f = 0;
+
+	char relativeFileName[1024];
+	{
+		const char* prefix[]={"./","../","../../","../../../","../../../../"};
+		int numPrefixes = sizeof(prefix)/sizeof(char*);
+
+		for (int i=0;i<numPrefixes;i++)
+		{
+			
+			sprintf(relativeFileName,"%s%s",prefix[i],fileName);
+			f = fopen(relativeFileName,"r");
+			if (f)
+			{
+				fclose(f);
+				break;
+			}
+		}
+	}
+
+	if (f)
+		fclose(f);
+	else
+		return;
+
+	objData->load(relativeFileName);
+
+	if (objData->objectCount>0)
+	{
+
+
+		int strideInBytes = 9*sizeof(float);
+
+		btAlignedObjectArray<GLInstanceVertex> vertexArray;
+		btAlignedObjectArray<int> indexArray;
+
+	
+
+
+		//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
+		int group=1;
+		int mask=1;
+		int index=0;
+		int colIndex = 0;
+
+		btAlignedObjectArray<GLInstanceVertex> vertices;
+		int stride2 = sizeof(GLInstanceVertex);
+		btAssert(stride2 == strideInBytes);
+
+		{
+		
+		
+			btAlignedObjectArray<btGpuChildShape> childShapes;
+
+			int numChildShapes = objData->objectCount;
+
+			for (int i=0;i<numChildShapes;i++)
+//			int i=4;
+			{
+
+				obj_object* object = objData->objectList[i];
+				int numVertices = i==numChildShapes-1? objData->vertexCount-object->vertex_offset : objData->objectList[i+1]->vertex_offset - object->vertex_offset;
+				int numFaces = i==numChildShapes-1? objData->faceCount - object->face_offset : objData->objectList[i+1]->face_offset-object->face_offset;
+
+
+
+
+				//for now, only support polyhedral child shapes
+				btGpuChildShape child;
+			
+				btVector3 pos(0,0,0);
+				btQuaternion orn(0,0,0,1);
+				for (int v=0;v<4;v++)
+				{
+					child.m_childPosition[v] = pos[v];
+					child.m_childOrientation[v] = orn[v];
+				}
+				
+				btTransform tr;
+				tr.setIdentity();
+				tr.setOrigin(pos);
+				tr.setRotation(orn);
+
+				int baseIndex = vertexArray.size();
+				
+				for (int f=0;f<numFaces;f++)
+				{
+					obj_face* face = objData->faceList[object->face_offset+f];
+					if (face->vertex_count==3)
+					{
+						for (int i=0;i<3;i++)
+						{
+							indexArray.push_back(face->vertex_index[i]);//-object->vertex_offset);
+						}
+					} else
+					{
+						btAssert(0);
+					}
+				}
+
+				btVector3 center(0,0,0);
+
+				btAlignedObjectArray<GLInstanceVertex> tmpVertices;
+				//add transformed graphics vertices and indices
+				btVector3 myScaling(1,1,1);//50,50,50);//300,300,300);
+				for (int v=0;v<numVertices;v++)
+				{
+					GLInstanceVertex vert;
+					obj_vector* orgVert = objData->vertexList[object->vertex_offset+v];
+					vert.uv[0] = 0.5f;
+					vert.uv[1] = 0.5f;
+					vert.normal[0]=0.f;
+					vert.normal[1]=1.f;
+					vert.normal[2]=0.f;
+					btVector3 vertPos;
+					vertPos[0] = orgVert->e[0]*myScaling[0];
+					vertPos[1] = orgVert->e[1]*myScaling[1];
+					vertPos[2] = orgVert->e[2]*myScaling[2];
+					vertPos[3] =0.f;
+					center+=vertPos;
+				}
+
+				center/=numVertices;
+
+				for (int v=0;v<numVertices;v++)
+				{
+					GLInstanceVertex vert;
+					obj_vector* orgVert = objData->vertexList[object->vertex_offset+v];
+					vert.uv[0] = 0.5f;
+					vert.uv[1] = 0.5f;
+					vert.normal[0]=0.f;
+					vert.normal[1]=1.f;
+					vert.normal[2]=0.f;
+					btVector3 vertPos;
+					vertPos[0] = orgVert->e[0]*myScaling[0];
+					vertPos[1] = orgVert->e[1]*myScaling[1];
+					vertPos[2] = orgVert->e[2]*myScaling[2];
+					vertPos[3] =0.f;
+	//				vertPos-=center;
+					vert.xyzw[0] = vertPos[0];
+					vert.xyzw[1] = vertPos[1];
+					vert.xyzw[2] = vertPos[2];
+
+					tmpVertices.push_back(vert);
+					btVector3 newPos = tr*vertPos;
+					vert.xyzw[0] = newPos[0];
+					vert.xyzw[1] = newPos[1];
+					vert.xyzw[2] = newPos[2];
+					vert.xyzw[3] = 0.f;
+					vertexArray.push_back(vert);
+				}
+
+
+				int childColIndex = m_data->m_np->registerConvexHullShape(&tmpVertices[0].xyzw[0],strideInBytes,numVertices, scaling);
+				child.m_shapeIndex = childColIndex;
+				childShapes.push_back(child);
+				colIndex = childColIndex;
+			}
+			//colIndex= m_data->m_np->registerCompoundShape(&childShapes);
+		
+		}
+
+		//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
+		int shapeId = ci.m_instancingRenderer->registerShape(&vertexArray[0].xyzw[0],vertexArray.size(),&indexArray[0],indexArray.size());
+
+		btVector4 colors[4] = 
+		{
+			btVector4(1,0,0,1),
+			btVector4(0,1,0,1),
+			btVector4(0,0,1,1),
+			btVector4(0,1,1,1),
+		};
+		
+		int curColor = 0;
+		for (int i=0;i<ci.arraySizeX;i++)
+		{
+			for (int j=0;j<ci.arraySizeY;j++)
+			{
+				for (int k=0;k<ci.arraySizeZ;k++)
+
+				{
+					float mass = 1;//j==0? 0.f : 1.f;
+
+					//btVector3 position(i*10*ci.gapX,j*ci.gapY,k*10*ci.gapZ);
+					btVector3 position(i*10*ci.gapX,1+j*ci.gapY,k*10*ci.gapZ);
+
+				//	btQuaternion orn(0,0,0,1);
+					btQuaternion orn(btVector3(0,0,1),1.3);
+				
+					btVector4 color = colors[curColor];
+					curColor++;
+					curColor&=3;
+					btVector4 scaling(1,1,1,1);
+					int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
+					int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(mass,position,orn,colIndex,index);
+				
+					index++;
+				}
+			}
+		}
+
+	}
+	delete objData;
+}
+
 void ConcaveCompoundScene::createDynamicObjects(const ConstructionInfo& ci)
 {
 	
diff --git a/demo/gpudemo/rigidbody/ConcaveScene.h b/demo/gpudemo/rigidbody/ConcaveScene.h
index 617cb5538..1fb7f0b6d 100644
--- a/demo/gpudemo/rigidbody/ConcaveScene.h
+++ b/demo/gpudemo/rigidbody/ConcaveScene.h
@@ -77,4 +77,25 @@ public:
 
 };
 
+
+class ConcaveCompound2Scene : public ConcaveCompoundScene
+{
+public:
+
+	ConcaveCompound2Scene(){}
+	virtual ~ConcaveCompound2Scene(){}
+	virtual const char* getName()
+	{
+		return "GRBConcave2Compound";
+	}
+
+	static GpuDemo* MyCreateFunc()
+	{
+		GpuDemo* demo = new ConcaveCompound2Scene;
+		return demo;
+	}
+	virtual void createDynamicObjects(const ConstructionInfo& ci);
+};
+
+
 #endif //CONCAVE_SCENE_H
diff --git a/opencl/gpu_sat/host/ConvexHullContact.cpp b/opencl/gpu_sat/host/ConvexHullContact.cpp
index 35a1bea9d..efc9e0598 100644
--- a/opencl/gpu_sat/host/ConvexHullContact.cpp
+++ b/opencl/gpu_sat/host/ConvexHullContact.cpp
@@ -905,14 +905,14 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 			hostCollidables[collidableIndexB].m_shapeType == SHAPE_CONVEX_HULL)
 		{
 			computeContactSphereConvex(i,bodyIndexA,bodyIndexB,collidableIndexA,collidableIndexB,&hostBodyBuf[0],
-				&hostCollidables[0],&hostConvexData[0],&hostVertices[0],&hostIndices[0],&hostFaces[0],&hostContacts[0],nContacts,nPairs);
+				&hostCollidables[0],&hostConvexData[0],&hostVertices[0],&hostIndices[0],&hostFaces[0],&hostContacts[0],nContacts,maxContactCapacity);
 		}
 
 		if (hostCollidables[collidableIndexA].m_shapeType == SHAPE_CONVEX_HULL &&
 			hostCollidables[collidableIndexB].m_shapeType == SHAPE_SPHERE)
 		{
 			computeContactSphereConvex(i,bodyIndexB,bodyIndexA,collidableIndexB,collidableIndexA,&hostBodyBuf[0],
-				&hostCollidables[0],&hostConvexData[0],&hostVertices[0],&hostIndices[0],&hostFaces[0],&hostContacts[0],nContacts,nPairs);
+				&hostCollidables[0],&hostConvexData[0],&hostVertices[0],&hostIndices[0],&hostFaces[0],&hostContacts[0],nContacts,maxContactCapacity);
 			//printf("convex-sphere\n");
 			
 		}
@@ -921,7 +921,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 			hostCollidables[collidableIndexB].m_shapeType == SHAPE_PLANE)
 		{
 			computeContactPlaneConvex(i,bodyIndexB,bodyIndexA,collidableIndexB,collidableIndexA,&hostBodyBuf[0],
-			&hostCollidables[0],&hostConvexData[0],&hostVertices[0],&hostIndices[0],&hostFaces[0],&hostContacts[0],nContacts,nPairs);
+			&hostCollidables[0],&hostConvexData[0],&hostVertices[0],&hostIndices[0],&hostFaces[0],&hostContacts[0],nContacts,maxContactCapacity);
 //			printf("convex-plane\n");
 			
 		}
@@ -930,7 +930,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 			hostCollidables[collidableIndexB].m_shapeType == SHAPE_CONVEX_HULL)
 		{
 			computeContactPlaneConvex(i,bodyIndexA,bodyIndexB,collidableIndexA,collidableIndexB,&hostBodyBuf[0],
-			&hostCollidables[0],&hostConvexData[0],&hostVertices[0],&hostIndices[0],&hostFaces[0],&hostContacts[0],nContacts,nPairs);
+			&hostCollidables[0],&hostConvexData[0],&hostVertices[0],&hostIndices[0],&hostFaces[0],&hostContacts[0],nContacts,maxContactCapacity);
 //			printf("plane-convex\n");
 			
 		}
@@ -940,7 +940,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 			hostCollidables[collidableIndexB].m_shapeType == SHAPE_PLANE)
 		{
 			computeContactPlaneCompound(i,bodyIndexB,bodyIndexA,collidableIndexB,collidableIndexA,&hostBodyBuf[0],
-			&hostCollidables[0],&hostConvexData[0],&hostVertices[0],&hostIndices[0],&hostFaces[0],&hostContacts[0],nContacts,nPairs);
+			&hostCollidables[0],&hostConvexData[0],&hostVertices[0],&hostIndices[0],&hostFaces[0],&hostContacts[0],nContacts,maxContactCapacity);
 //			printf("convex-plane\n");
 			
 		}
@@ -949,7 +949,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 			hostCollidables[collidableIndexB].m_shapeType == SHAPE_COMPOUND_OF_CONVEX_HULLS)
 		{
 			computeContactPlaneCompound(i,bodyIndexA,bodyIndexB,collidableIndexA,collidableIndexB,&hostBodyBuf[0],
-			&hostCollidables[0],&hostConvexData[0],&hostVertices[0],&hostIndices[0],&hostFaces[0],&hostContacts[0],nContacts,nPairs);
+			&hostCollidables[0],&hostConvexData[0],&hostVertices[0],&hostIndices[0],&hostFaces[0],&hostContacts[0],nContacts,maxContactCapacity);
 //			printf("plane-convex\n");
 			
 		}
@@ -1131,6 +1131,14 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 							int num = numConcavePairs;
 							launcher.launch1D( num);
 							clFinish(m_queue);
+
+							btAlignedObjectArray<btVector3> cpuCompoundSepNormals;
+							concaveSepNormals.copyToHost(cpuCompoundSepNormals);
+							btAlignedObjectArray<btInt4> cpuConcavePairs;
+							triangleConvexPairsOut.copyToHost(cpuConcavePairs);
+
+							printf("!\n");
+
 						}
 					}
 				}
@@ -1238,6 +1246,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 				int num = numCompoundPairs;
 				launcher.launch1D( num);
 				clFinish(m_queue);
+
 			
 			}
 
@@ -1261,7 +1270,7 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 			if (numConcavePairs)
 		{
 			BT_PROFILE("findConcaveSphereContactsKernel");
-
+				nContacts = m_totalContactsOut.at(0);
 			btBufferInfoCL bInfo[] = { 
 				btBufferInfoCL( triangleConvexPairsOut.getBufferCL() ), 
 				btBufferInfoCL( bodyBuf->getBufferCL(),true), 
@@ -1541,6 +1550,8 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 			btLauncherCL launcher(m_queue, m_clipCompoundsHullHullKernel);
 			launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(btBufferInfoCL) );
 			launcher.setConst( nCompoundsPairs  );
+			launcher.setConst(maxContactCapacity);
+
 			int num = nCompoundsPairs;
 			launcher.launch1D( num);
 			clFinish(m_queue);
diff --git a/opencl/gpu_sat/kernels/satClipHullContacts.cl b/opencl/gpu_sat/kernels/satClipHullContacts.cl
index 4d3f5140e..96fbdfa25 100644
--- a/opencl/gpu_sat/kernels/satClipHullContacts.cl
+++ b/opencl/gpu_sat/kernels/satClipHullContacts.cl
@@ -1062,7 +1062,7 @@ __kernel void   clipCompoundsHullHullKernel( __global const int4* gpuCompoundPai
 																					__global const int* gpuHasCompoundSepNormalsOut,
 																					__global Contact4* restrict globalContactsOut,
 																					counter32_t nGlobalContactsOut,
-																					int numCompoundPairs)
+																					int numCompoundPairs, int maxContactCapacity)
 {
 
 	int i = get_global_id(0);
@@ -1155,7 +1155,7 @@ __kernel void   clipCompoundsHullHullKernel( __global const int4* gpuCompoundPai
 		
 				int dstIdx;
 				AppendInc( nGlobalContactsOut, dstIdx );
-				//if ((dstIdx+nReducedContacts) < capacity)
+				if ((dstIdx+nReducedContacts) < maxContactCapacity)
 				{
 					__global Contact4* c = globalContactsOut+ dstIdx;
 					c->m_worldNormal = normal;
diff --git a/opencl/gpu_sat/kernels/satClipHullContacts.h b/opencl/gpu_sat/kernels/satClipHullContacts.h
index c8be64714..4dc53cf8d 100644
--- a/opencl/gpu_sat/kernels/satClipHullContacts.h
+++ b/opencl/gpu_sat/kernels/satClipHullContacts.h
@@ -1064,7 +1064,7 @@ static const char* satClipKernelsCL= \
 "																					__global const int* gpuHasCompoundSepNormalsOut,\n"
 "																					__global Contact4* restrict globalContactsOut,\n"
 "																					counter32_t nGlobalContactsOut,\n"
-"																					int numCompoundPairs)\n"
+"																					int numCompoundPairs, int maxContactCapacity)\n"
 "{\n"
 "\n"
 "	int i = get_global_id(0);\n"
@@ -1157,7 +1157,7 @@ static const char* satClipKernelsCL= \
 "		\n"
 "				int dstIdx;\n"
 "				AppendInc( nGlobalContactsOut, dstIdx );\n"
-"				//if ((dstIdx+nReducedContacts) < capacity)\n"
+"				if ((dstIdx+nReducedContacts) < maxContactCapacity)\n"
 "				{\n"
 "					__global Contact4* c = globalContactsOut+ dstIdx;\n"
 "					c->m_worldNormal = normal;\n"