allow multiple concave meshes, each with their own bvh

demo/gpudemo/main_opengl3core.cpp

@@ -66,7 +66,7 @@ int selectedDemo = 0;
 GpuDemo::CreateFunc* allDemos[]=
 {
 //		ConcaveCompound2Scene::MyCreateFunc,
-
+		ConcaveSphereScene::MyCreateFunc,
 
 
 	GpuBoxPlaneScene::MyCreateFunc,	

demo/gpudemo/rigidbody/ConcaveScene.cpp

@@ -155,21 +155,10 @@ GLInstanceGraphicsShape* createGraphicsShapeFromWavefrontObj(objLoader* obj)
 }
 
 
-void ConcaveScene::createConcaveMesh(const ConstructionInfo& ci)
+void ConcaveScene::createConcaveMesh(const ConstructionInfo& ci, const char* fileName, const b3Vector3& shift, const b3Vector3& scaling)
 {
 	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/teddy2_VHACD_CHs.obj";
 	
-	b3Vector3 shift(0,0,0);//0,230,80);//150,-100,-120);
-	btVector4 scaling(4,4,4,1);
 	FILE* f = 0;
 
 	char relativeFileName[1024];
@@ -254,7 +243,27 @@ void ConcaveScene::setupScene(const ConstructionInfo& ci)
 
 	if (1)
 	{
-		createConcaveMesh(ci);
+
+		//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/teddy2_VHACD_CHs.obj";
+	
+		b3Vector3 shift1(0,-50,0);//0,230,80);//150,-100,-120);
+		
+		btVector4 scaling(4,4,4,1);
+
+		createConcaveMesh(ci,"data/plane100.obj",shift1,scaling);
+		//createConcaveMesh(ci,"data/plane100.obj",shift,scaling);
+
+		b3Vector3 shift2(0,0,0);//0,230,80);//150,-100,-120);
+		createConcaveMesh(ci,"data/teddy.obj",shift2,scaling);
+				
 	} else
 	{
 		int strideInBytes = 9*sizeof(float);

demo/gpudemo/rigidbody/ConcaveScene.h

@@ -2,6 +2,7 @@
 #define CONCAVE_SCENE_H
 
 #include "GpuRigidBodyDemo.h"
+#include "Bullet3Common/b3Vector3.h"
 
 class ConcaveScene : public GpuRigidBodyDemo
 {
@@ -24,8 +25,7 @@ public:
 
 	virtual void createDynamicObjects(const ConstructionInfo& ci);
 
-	virtual void createConcaveMesh(const ConstructionInfo& ci);
+	virtual void createConcaveMesh(const ConstructionInfo& ci, const char* fileName, const b3Vector3& shift, const b3Vector3& scaling);
-
 	
 };
 

opencl/gpu_narrowphase/host/b3BvhInfo.h

@@ -0,0 +1,18 @@
+#ifndef B3_BVH_INFO_H
+#define B3_BVH_INFO_H
+
+#include "Bullet3Common/b3Vector3.h"
+
+struct b3BvhInfo
+{
+	b3Vector3	m_aabbMin;
+	b3Vector3	m_aabbMax;
+	b3Vector3	m_quantization;
+	int			m_numNodes;
+	int			m_numSubTrees;
+	int			m_nodeOffset;
+	int			m_subTreeOffset;
+
+};
+
+#endif //B3_BVH_INFO_H

opencl/gpu_narrowphase/host/b3Collidable.h

@@ -16,7 +16,10 @@ enum btShapeTypes
 
 struct b3Collidable
 {
-	int m_numChildShapes;
+	union {
+		int m_numChildShapes;
+		int m_bvhIndex;
+	};
 	float m_radius;
 	int m_shapeType;
 	int m_shapeIndex;

opencl/gpu_narrowphase/host/b3ConvexHullContact.cpp

@@ -38,6 +38,7 @@ typedef b3AlignedObjectArray<b3Vector3> btVertexArray;
 #include "../kernels/bvhTraversal.h"
 #include "../kernels/primitiveContacts.h"
 
+
 #include "Bullet3Geometry/b3AabbUtil.h"
 
 
@@ -853,6 +854,8 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 			b3AlignedObjectArray<class b3OptimizedBvh*>& bvhData,
 			btOpenCLArray<btQuantizedBvhNode>*	treeNodesGPU,
 			btOpenCLArray<btBvhSubtreeInfo>*	subTreesGPU,
+			btOpenCLArray<b3BvhInfo>*	bvhInfo,
+
 			int numObjects,
 			int maxTriConvexPairCapacity,
 			btOpenCLArray<btInt4>& triangleConvexPairsOut,
@@ -1077,13 +1080,11 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 			//now perform the tree query on GPU
 			{
 				
-				int numNodes = bvhData.size()? bvhData[0]->getQuantizedNodeArray().size() : 0;
+				
-				if (numNodes)
+				
 				{
-					int numSubTrees = subTreesGPU->size();
+				
-					b3Vector3 bvhAabbMin = bvhData[0]->m_bvhAabbMin;
+
-					b3Vector3 bvhAabbMax = bvhData[0]->m_bvhAabbMax;
-					b3Vector3 bvhQuantization = bvhData[0]->m_bvhQuantization;
 					{
 						BT_PROFILE("m_bvhTraversalKernel");
 						numConcavePairs = numConcavePairsOut.at(0);
@@ -1096,10 +1097,8 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 						launcher.setBuffer( numConcavePairsOut.getBufferCL());
 						launcher.setBuffer( subTreesGPU->getBufferCL());
 						launcher.setBuffer( treeNodesGPU->getBufferCL());
-						launcher.setConst( bvhAabbMin);
+						launcher.setBuffer( bvhInfo->getBufferCL());
-						launcher.setConst( bvhAabbMax);
+						
-						launcher.setConst( bvhQuantization);
-						launcher.setConst(numSubTrees);
 						launcher.setConst( nPairs  );
 						launcher.setConst( maxTriConvexPairCapacity);
 						int num = nPairs;

opencl/gpu_narrowphase/host/b3ConvexHullContact.h

@@ -12,6 +12,7 @@
 #include "parallel_primitives/host/btInt2.h"
 #include "parallel_primitives/host/btInt4.h"
 #include "b3OptimizedBvh.h"
+#include "b3BvhInfo.h"
 
 //#include "../../dynamics/basic_demo/Stubs/ChNarrowPhase.h"
 
@@ -85,6 +86,7 @@ struct GpuSatCollision
 		   b3AlignedObjectArray<class b3OptimizedBvh*>& bvhData,
 		   btOpenCLArray<btQuantizedBvhNode>*	treeNodesGPU,
 			btOpenCLArray<btBvhSubtreeInfo>*	subTreesGPU,
+			btOpenCLArray<b3BvhInfo>*	bvhInfo,
 			int numObjects,
 			int maxTriConvexPairCapacity,
 			btOpenCLArray<btInt4>& triangleConvexPairs,

opencl/gpu_narrowphase/kernels/bvhTraversal.cl

@@ -21,6 +21,19 @@ typedef struct
 	//4 bytes
 	int	m_escapeIndexOrTriangleIndex;
 } btQuantizedBvhNode;
+
+typedef struct
+{
+	float4		m_aabbMin;
+	float4		m_aabbMax;
+	float4		m_quantization;
+	int			m_numNodes;
+	int			m_numSubTrees;
+	int			m_nodeOffset;
+	int			m_subTreeOffset;
+
+} b3BvhInfo;
+
 /*
 	bool isLeafNode() const
 	{
@@ -185,12 +198,9 @@ __kernel void   bvhTraversalKernel( __global const int2* pairs,
 									__global btAabbCL* aabbs,
 									__global int4* concavePairsOut,
 									__global volatile int* numConcavePairsOut,
-									__global const btBvhSubtreeInfo* subtreeHeaders,
+									__global const btBvhSubtreeInfo* subtreeHeadersRoot,
-									__global const btQuantizedBvhNode* quantizedNodes,
+									__global const btQuantizedBvhNode* quantizedNodesRoot,
-									float4 bvhAabbMin,
+									__global const b3BvhInfo* bvhInfos,
-									float4 bvhAabbMax,
-									float4 bvhQuantization,
-									int numSubtreeHeaders,
 									int numPairs,
 									int maxNumConcavePairsCapacity)
 {
@@ -220,7 +230,16 @@ __kernel void   bvhTraversalKernel( __global const int2* pairs,
 		)
 		return;
 
+	b3BvhInfo bvhInfo = bvhInfos[collidables[collidableIndexA].m_numChildShapes];
+
+	float4 bvhAabbMin = bvhInfo.m_aabbMin;
+	float4 bvhAabbMax = bvhInfo.m_aabbMax;
+	float4 bvhQuantization = bvhInfo.m_quantization;
+	int numSubtreeHeaders = bvhInfo.m_numSubTrees;
+	__global const btBvhSubtreeInfo* subtreeHeaders = &subtreeHeadersRoot[bvhInfo.m_subTreeOffset];
+	__global const btQuantizedBvhNode* quantizedNodes = &quantizedNodesRoot[bvhInfo.m_nodeOffset];
 	
+
 	unsigned short int quantizedQueryAabbMin[3];
 	unsigned short int quantizedQueryAabbMax[3];
 	quantizeWithClamp(quantizedQueryAabbMin,aabbs[bodyIndexB].m_min,false,bvhAabbMin, bvhAabbMax,bvhQuantization);

opencl/gpu_narrowphase/kernels/bvhTraversal.h

@@ -23,6 +23,19 @@ static const char* bvhTraversalKernelCL= \
 "	//4 bytes\n"
 "	int	m_escapeIndexOrTriangleIndex;\n"
 "} btQuantizedBvhNode;\n"
+"\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"
+"\n"
+"} b3BvhInfo;\n"
+"\n"
 "/*\n"
 "	bool isLeafNode() const\n"
 "	{\n"
@@ -187,12 +200,9 @@ static const char* bvhTraversalKernelCL= \
 "									__global btAabbCL* aabbs,\n"
 "									__global int4* concavePairsOut,\n"
 "									__global volatile int* numConcavePairsOut,\n"
-"									__global const btBvhSubtreeInfo* subtreeHeaders,\n"
+"									__global const btBvhSubtreeInfo* subtreeHeadersRoot,\n"
-"									__global const btQuantizedBvhNode* quantizedNodes,\n"
+"									__global const btQuantizedBvhNode* quantizedNodesRoot,\n"
-"									float4 bvhAabbMin,\n"
+"									__global const b3BvhInfo* bvhInfos,\n"
-"									float4 bvhAabbMax,\n"
-"									float4 bvhQuantization,\n"
-"									int numSubtreeHeaders,\n"
 "									int numPairs,\n"
 "									int maxNumConcavePairsCapacity)\n"
 "{\n"
@@ -222,7 +232,16 @@ static const char* bvhTraversalKernelCL= \
 "		)\n"
 "		return;\n"
 "\n"
+"	b3BvhInfo bvhInfo = bvhInfos[collidables[collidableIndexA].m_numChildShapes];\n"
+"\n"
+"	float4 bvhAabbMin = bvhInfo.m_aabbMin;\n"
+"	float4 bvhAabbMax = bvhInfo.m_aabbMax;\n"
+"	float4 bvhQuantization = bvhInfo.m_quantization;\n"
+"	int numSubtreeHeaders = bvhInfo.m_numSubTrees;\n"
+"	__global const btBvhSubtreeInfo* subtreeHeaders = &subtreeHeadersRoot[bvhInfo.m_subTreeOffset];\n"
+"	__global const btQuantizedBvhNode* quantizedNodes = &quantizedNodesRoot[bvhInfo.m_nodeOffset];\n"
 "	\n"
+"\n"
 "	unsigned short int quantizedQueryAabbMin[3];\n"
 "	unsigned short int quantizedQueryAabbMax[3];\n"
 "	quantizeWithClamp(quantizedQueryAabbMin,aabbs[bodyIndexB].m_min,false,bvhAabbMin, bvhAabbMax,bvhQuantization);\n"

opencl/gpu_rigidbody/host/b3GpuNarrowPhase.cpp

@@ -10,6 +10,7 @@
 #include "../../gpu_narrowphase/host/b3OptimizedBvh.h"
 #include "../../gpu_narrowphase/host/b3TriangleIndexVertexArray.h"
 #include "Bullet3Geometry/b3AabbUtil.h"
+#include "../../gpu_narrowphase/host/b3BvhInfo.h"
 
 struct btGpuNarrowPhaseInternalData
 {
@@ -64,6 +65,13 @@ struct btGpuNarrowPhaseInternalData
 	b3AlignedObjectArray<b3SapAabb>* m_localShapeAABBCPU;
 
 	b3AlignedObjectArray<class b3OptimizedBvh*> m_bvhData;
+
+	b3AlignedObjectArray<btQuantizedBvhNode>	m_treeNodesCPU;
+	b3AlignedObjectArray<btBvhSubtreeInfo>	m_subTreesCPU;
+
+	b3AlignedObjectArray<b3BvhInfo>	m_bvhInfoCPU;
+	btOpenCLArray<b3BvhInfo>*			m_bvhInfoGPU;
+	
 	btOpenCLArray<btQuantizedBvhNode>*	m_treeNodesGPU;
 	btOpenCLArray<btBvhSubtreeInfo>*	m_subTreesGPU;
 	
@@ -141,9 +149,10 @@ m_queue(queue)
 	m_data->m_numAcceleratedShapes = 0;
 	m_data->m_numAcceleratedRigidBodies = 0;
     
-	m_data->m_treeNodesGPU = 0;
+		
-	m_data->m_subTreesGPU = 0;
+	m_data->m_subTreesGPU = new btOpenCLArray<btBvhSubtreeInfo>(this->m_context,this->m_queue);
-	
+	m_data->m_treeNodesGPU = new btOpenCLArray<btQuantizedBvhNode>(this->m_context,this->m_queue);
+	m_data->m_bvhInfoGPU = new btOpenCLArray<b3BvhInfo>(this->m_context,this->m_queue);
 
 	//m_data->m_contactCGPU = new btOpenCLArray<Constraint4>(ctx,queue,config.m_maxBroadphasePairs,false);
 	//m_data->m_frictionCGPU = new btOpenCLArray<adl::Solver<adl::TYPE_CL>::allocateFrictionConstraint( m_data->m_deviceCL, config.m_maxBroadphasePairs);
@@ -178,8 +187,11 @@ b3GpuNarrowPhase::~b3GpuNarrowPhase()
 	delete m_data->m_worldVertsA1GPU;
     delete m_data->m_worldVertsB2GPU;
     
+	delete m_data->m_bvhInfoGPU;
+
 	delete m_data->m_treeNodesGPU;
 	delete m_data->m_subTreesGPU;
+
     
     delete m_data->m_convexData;
 	delete m_data;
@@ -487,14 +499,7 @@ int		b3GpuNarrowPhase::registerCompoundShape(b3AlignedObjectArray<btGpuChildShap
 
 int		b3GpuNarrowPhase::registerConcaveMesh(b3AlignedObjectArray<b3Vector3>* vertices, b3AlignedObjectArray<int>* indices,const float* scaling1)
 {
-	//right now we only support one single mesh, it is on the todo to merge all mesh data etc
+	
-	btAssert(m_data->m_treeNodesGPU ==0);
-	btAssert(m_data->m_subTreesGPU ==0);
-	if (m_data->m_treeNodesGPU)
-	{
-		printf("error, only 1 single concave mesh supported at the moment\n");
-		exit (0);
-	}
 
 	b3Vector3 scaling(scaling1[0],scaling1[1],scaling1[2]);
 
@@ -503,8 +508,8 @@ int		b3GpuNarrowPhase::registerConcaveMesh(b3AlignedObjectArray<b3Vector3>* vert
 	
 	col.m_shapeType = SHAPE_CONCAVE_TRIMESH;
 	col.m_shapeIndex = registerConcaveMeshShape(vertices,indices,col,scaling);
-
+	col.m_bvhIndex = m_data->m_bvhInfoCPU.size();
-	
+		
 
 	b3SapAabb aabb;
 	b3Vector3 myAabbMin(1e30f,1e30f,1e30f);
@@ -546,15 +551,42 @@ int		b3GpuNarrowPhase::registerConcaveMesh(b3AlignedObjectArray<b3Vector3>* vert
 	bvh->build(meshInterface, useQuantizedAabbCompression, (b3Vector3&)aabb.m_min, (b3Vector3&)aabb.m_max);
 	m_data->m_bvhData.push_back(bvh);
 	int numNodes = bvh->getQuantizedNodeArray().size();
-	btOpenCLArray<btQuantizedBvhNode>*	treeNodesGPU = new btOpenCLArray<btQuantizedBvhNode>(this->m_context,this->m_queue,numNodes);
+	//btOpenCLArray<btQuantizedBvhNode>*	treeNodesGPU = new btOpenCLArray<btQuantizedBvhNode>(this->m_context,this->m_queue,numNodes);
-	treeNodesGPU->copyFromHost(bvh->getQuantizedNodeArray());
+	//treeNodesGPU->copyFromHost(bvh->getQuantizedNodeArray());
-
 	int numSubTrees = bvh->getSubtreeInfoArray().size();
-	btOpenCLArray<btBvhSubtreeInfo>* subTreesGPU = new btOpenCLArray<btBvhSubtreeInfo>(this->m_context,this->m_queue,numSubTrees);
-	subTreesGPU->copyFromHost(bvh->getSubtreeInfoArray());
 
-	m_data->m_treeNodesGPU = treeNodesGPU;
+	b3BvhInfo bvhInfo;
-	m_data->m_subTreesGPU = subTreesGPU;
+	
+	bvhInfo.m_aabbMin = bvh->m_bvhAabbMin;
+	bvhInfo.m_aabbMax = bvh->m_bvhAabbMax;
+	bvhInfo.m_quantization = bvh->m_bvhQuantization;
+	bvhInfo.m_numNodes = numNodes;
+	bvhInfo.m_numSubTrees = numSubTrees;
+	bvhInfo.m_nodeOffset = m_data->m_treeNodesCPU.size();
+	bvhInfo.m_subTreeOffset = m_data->m_subTreesCPU.size();
+
+	m_data->m_bvhInfoCPU.push_back(bvhInfo);
+	m_data->m_bvhInfoGPU->copyFromHost(m_data->m_bvhInfoCPU);
+
+
+	int numNewSubtrees = bvh->getSubtreeInfoArray().size();
+	m_data->m_subTreesCPU.reserve(m_data->m_subTreesCPU.size()+numNewSubtrees);
+	for (int i=0;i<numNewSubtrees;i++)
+	{
+		m_data->m_subTreesCPU.push_back(bvh->getSubtreeInfoArray()[i]);
+	}
+	int numNewTreeNodes = bvh->getQuantizedNodeArray().size();
+
+	for (int i=0;i<numNewTreeNodes;i++)
+	{
+		m_data->m_treeNodesCPU.push_back(bvh->getQuantizedNodeArray()[i]);
+	}
+
+	//btOpenCLArray<btBvhSubtreeInfo>* subTreesGPU = new btOpenCLArray<btBvhSubtreeInfo>(this->m_context,this->m_queue,numSubTrees);
+	//subTreesGPU->copyFromHost(bvh->getSubtreeInfoArray());
+
+	m_data->m_treeNodesGPU->copyFromHost(m_data->m_treeNodesCPU);
+	m_data->m_subTreesGPU->copyFromHost(m_data->m_subTreesCPU);
 
 
 	return collidableIndex;
@@ -739,6 +771,7 @@ void b3GpuNarrowPhase::computeContacts(cl_mem broadphasePairs, int numBroadphase
 		m_data->m_bvhData,
 		m_data->m_treeNodesGPU,
 		m_data->m_subTreesGPU,
+		m_data->m_bvhInfoGPU,
 		numObjects,
 		maxTriConvexPairCapacity,
 		triangleConvexPairs,