Bart/bullet3
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

extend Wavefront obj loader to recognize objects (for compound shape creation)

Browse Source 
add reproduction of a bug in compound versus concave trimesh
This commit is contained in:
erwin coumans
2013-04-12 10:25:11 -07:00
parent 467a68293b
commit ec0d2ed523
13 changed files with 386 additions and 21 deletions
Download Patch File Download Diff File Expand all files Collapse all files

11
data/cube.mtl Normal file
View File

@@ -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

26
data/cube.obj Normal file
View File

@@ -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

3
demo/ObjLoader/objLoader.cpp
View File

@@ -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;

2
demo/ObjLoader/objLoader.h
View File

@@ -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;

29
demo/ObjLoader/obj_parser.cpp
View File

@@ -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;

16
demo/ObjLoader/obj_parser.h
View File

@@ -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;

6
demo/gpudemo/GpuDemo.h
View File

@@ -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),

11
demo/gpudemo/main_opengl3core.cpp
View File

@@ -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,

249
demo/gpudemo/rigidbody/ConcaveScene.cpp
View File

@@ -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)
{

21
demo/gpudemo/rigidbody/ConcaveScene.h
View File

@@ -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

25
opencl/gpu_sat/host/ConvexHullContact.cpp
View File

@@ -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);

4
opencl/gpu_sat/kernels/satClipHullContacts.cl
View File

@@ -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;

4
opencl/gpu_sat/kernels/satClipHullContacts.h
View File

@@ -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"