From fbd192f360a8db4450ccce2e522f6aa01892b6c1 Mon Sep 17 00:00:00 2001 From: erwin coumans Date: Fri, 19 Apr 2013 23:30:29 -0700 Subject: [PATCH] allow multiple concave meshes, each with their own bvh --- demo/gpudemo/main_opengl3core.cpp | 2 +- demo/gpudemo/rigidbody/ConcaveScene.cpp | 35 +++++---- demo/gpudemo/rigidbody/ConcaveScene.h | 4 +- opencl/gpu_narrowphase/host/b3BvhInfo.h | 18 +++++ opencl/gpu_narrowphase/host/b3Collidable.h | 5 +- .../host/b3ConvexHullContact.cpp | 19 +++-- .../host/b3ConvexHullContact.h | 2 + .../gpu_narrowphase/kernels/bvhTraversal.cl | 31 ++++++-- opencl/gpu_narrowphase/kernels/bvhTraversal.h | 31 ++++++-- .../gpu_rigidbody/host/b3GpuNarrowPhase.cpp | 73 ++++++++++++++----- 10 files changed, 161 insertions(+), 59 deletions(-) create mode 100644 opencl/gpu_narrowphase/host/b3BvhInfo.h diff --git a/demo/gpudemo/main_opengl3core.cpp b/demo/gpudemo/main_opengl3core.cpp index ec97fccc1..edacb04bf 100644 --- a/demo/gpudemo/main_opengl3core.cpp +++ b/demo/gpudemo/main_opengl3core.cpp @@ -66,7 +66,7 @@ int selectedDemo = 0; GpuDemo::CreateFunc* allDemos[]= { // ConcaveCompound2Scene::MyCreateFunc, - + ConcaveSphereScene::MyCreateFunc, GpuBoxPlaneScene::MyCreateFunc, diff --git a/demo/gpudemo/rigidbody/ConcaveScene.cpp b/demo/gpudemo/rigidbody/ConcaveScene.cpp index a077306db..e401020f4 100644 --- a/demo/gpudemo/rigidbody/ConcaveScene.cpp +++ b/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); diff --git a/demo/gpudemo/rigidbody/ConcaveScene.h b/demo/gpudemo/rigidbody/ConcaveScene.h index 1fb7f0b6d..47fe8597b 100644 --- a/demo/gpudemo/rigidbody/ConcaveScene.h +++ b/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); }; diff --git a/opencl/gpu_narrowphase/host/b3BvhInfo.h b/opencl/gpu_narrowphase/host/b3BvhInfo.h new file mode 100644 index 000000000..872f03950 --- /dev/null +++ b/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 \ No newline at end of file diff --git a/opencl/gpu_narrowphase/host/b3Collidable.h b/opencl/gpu_narrowphase/host/b3Collidable.h index 69805617d..a483abb92 100644 --- a/opencl/gpu_narrowphase/host/b3Collidable.h +++ b/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; diff --git a/opencl/gpu_narrowphase/host/b3ConvexHullContact.cpp b/opencl/gpu_narrowphase/host/b3ConvexHullContact.cpp index 4b931b5c3..c8e63b676 100644 --- a/opencl/gpu_narrowphase/host/b3ConvexHullContact.cpp +++ b/opencl/gpu_narrowphase/host/b3ConvexHullContact.cpp @@ -38,6 +38,7 @@ typedef b3AlignedObjectArray btVertexArray; #include "../kernels/bvhTraversal.h" #include "../kernels/primitiveContacts.h" + #include "Bullet3Geometry/b3AabbUtil.h" @@ -853,6 +854,8 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray& bvhData, btOpenCLArray* treeNodesGPU, btOpenCLArray* subTreesGPU, + btOpenCLArray* bvhInfo, + int numObjects, int maxTriConvexPairCapacity, btOpenCLArray& triangleConvexPairsOut, @@ -1077,13 +1080,11 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArraygetQuantizedNodeArray().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 btOpenCLArraygetBufferCL()); launcher.setBuffer( treeNodesGPU->getBufferCL()); - launcher.setConst( bvhAabbMin); - launcher.setConst( bvhAabbMax); - launcher.setConst( bvhQuantization); - launcher.setConst(numSubTrees); + launcher.setBuffer( bvhInfo->getBufferCL()); + launcher.setConst( nPairs ); launcher.setConst( maxTriConvexPairCapacity); int num = nPairs; diff --git a/opencl/gpu_narrowphase/host/b3ConvexHullContact.h b/opencl/gpu_narrowphase/host/b3ConvexHullContact.h index bda0fcbda..bfdf955e6 100644 --- a/opencl/gpu_narrowphase/host/b3ConvexHullContact.h +++ b/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& bvhData, btOpenCLArray* treeNodesGPU, btOpenCLArray* subTreesGPU, + btOpenCLArray* bvhInfo, int numObjects, int maxTriConvexPairCapacity, btOpenCLArray& triangleConvexPairs, diff --git a/opencl/gpu_narrowphase/kernels/bvhTraversal.cl b/opencl/gpu_narrowphase/kernels/bvhTraversal.cl index 33795628a..db6257f15 100644 --- a/opencl/gpu_narrowphase/kernels/bvhTraversal.cl +++ b/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 btQuantizedBvhNode* quantizedNodes, - float4 bvhAabbMin, - float4 bvhAabbMax, - float4 bvhQuantization, - int numSubtreeHeaders, + __global const btBvhSubtreeInfo* subtreeHeadersRoot, + __global const btQuantizedBvhNode* quantizedNodesRoot, + __global const b3BvhInfo* bvhInfos, 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); diff --git a/opencl/gpu_narrowphase/kernels/bvhTraversal.h b/opencl/gpu_narrowphase/kernels/bvhTraversal.h index 93fb018e4..6640917cb 100644 --- a/opencl/gpu_narrowphase/kernels/bvhTraversal.h +++ b/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 btQuantizedBvhNode* quantizedNodes,\n" -" float4 bvhAabbMin,\n" -" float4 bvhAabbMax,\n" -" float4 bvhQuantization,\n" -" int numSubtreeHeaders,\n" +" __global const btBvhSubtreeInfo* subtreeHeadersRoot,\n" +" __global const btQuantizedBvhNode* quantizedNodesRoot,\n" +" __global const b3BvhInfo* bvhInfos,\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" diff --git a/opencl/gpu_rigidbody/host/b3GpuNarrowPhase.cpp b/opencl/gpu_rigidbody/host/b3GpuNarrowPhase.cpp index 0722e6126..448d22d47 100644 --- a/opencl/gpu_rigidbody/host/b3GpuNarrowPhase.cpp +++ b/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* m_localShapeAABBCPU; b3AlignedObjectArray m_bvhData; + + b3AlignedObjectArray m_treeNodesCPU; + b3AlignedObjectArray m_subTreesCPU; + + b3AlignedObjectArray m_bvhInfoCPU; + btOpenCLArray* m_bvhInfoGPU; + btOpenCLArray* m_treeNodesGPU; btOpenCLArray* 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(this->m_context,this->m_queue); + m_data->m_treeNodesGPU = new btOpenCLArray(this->m_context,this->m_queue); + m_data->m_bvhInfoGPU = new btOpenCLArray(this->m_context,this->m_queue); //m_data->m_contactCGPU = new btOpenCLArray(ctx,queue,config.m_maxBroadphasePairs,false); //m_data->m_frictionCGPU = new btOpenCLArray::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* vertices, b3AlignedObjectArray* 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* 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* 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* treeNodesGPU = new btOpenCLArray(this->m_context,this->m_queue,numNodes); - treeNodesGPU->copyFromHost(bvh->getQuantizedNodeArray()); - + //btOpenCLArray* treeNodesGPU = new btOpenCLArray(this->m_context,this->m_queue,numNodes); + //treeNodesGPU->copyFromHost(bvh->getQuantizedNodeArray()); int numSubTrees = bvh->getSubtreeInfoArray().size(); - btOpenCLArray* subTreesGPU = new btOpenCLArray(this->m_context,this->m_queue,numSubTrees); - subTreesGPU->copyFromHost(bvh->getSubtreeInfoArray()); - m_data->m_treeNodesGPU = treeNodesGPU; - m_data->m_subTreesGPU = subTreesGPU; + b3BvhInfo bvhInfo; + + 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;im_subTreesCPU.push_back(bvh->getSubtreeInfoArray()[i]); + } + int numNewTreeNodes = bvh->getQuantizedNodeArray().size(); + + for (int i=0;im_treeNodesCPU.push_back(bvh->getQuantizedNodeArray()[i]); + } + + //btOpenCLArray* subTreesGPU = new btOpenCLArray(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,