From ad8585f1841ad9ea334ee3311a56f0d1f4a723ae Mon Sep 17 00:00:00 2001 From: erwincoumans Date: Sun, 10 Nov 2013 14:15:23 -0800 Subject: [PATCH] prepare for btGpuGridBroadphase --- Demos3/GpuDemos/GpuDemo.h | 6 +- Demos3/GpuDemos/broadphase/PairBench.cpp | 20 ++-- Demos3/GpuDemos/main_opengl3core.cpp | 2 +- .../GpuDemos/rigidbody/GpuRigidBodyDemo.cpp | 5 + .../rigidbody/GpuRigidBodyDemoInternalData.h | 2 +- Demos3/ImplicitCloth/premake4.lua | 5 +- Demos3/SimpleOpenGL3/premake4.lua | 5 +- .../b3GpuBroadphaseInterface.h | 39 ++++++++ .../b3GpuGridBroadphase.cpp | 98 +++++++++++++++++++ .../BroadphaseCollision/b3GpuGridBroadphase.h | 44 +++++++++ .../BroadphaseCollision/b3GpuSapBroadphase.h | 30 ++++-- .../RigidBody/b3GpuRigidBodyPipeline.cpp | 17 ++-- .../RigidBody/b3GpuRigidBodyPipeline.h | 2 +- .../b3GpuRigidBodyPipelineInternalData.h | 2 +- 14 files changed, 241 insertions(+), 36 deletions(-) create mode 100644 src/Bullet3OpenCL/BroadphaseCollision/b3GpuBroadphaseInterface.h create mode 100644 src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.cpp create mode 100644 src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.h diff --git a/Demos3/GpuDemos/GpuDemo.h b/Demos3/GpuDemos/GpuDemo.h index 0cfe70ef0..f779f34cb 100644 --- a/Demos3/GpuDemos/GpuDemo.h +++ b/Demos3/GpuDemos/GpuDemo.h @@ -48,9 +48,9 @@ public: arraySizeZ(10), #else - arraySizeX(30), - arraySizeY(30), - arraySizeZ(30), + arraySizeX(1), + arraySizeY(1), + arraySizeZ(1), #endif m_useConcaveMesh(false), gapX(16.3), diff --git a/Demos3/GpuDemos/broadphase/PairBench.cpp b/Demos3/GpuDemos/broadphase/PairBench.cpp index a31b46002..2fe4ab5ad 100644 --- a/Demos3/GpuDemos/broadphase/PairBench.cpp +++ b/Demos3/GpuDemos/broadphase/PairBench.cpp @@ -4,6 +4,8 @@ #include "Bullet3Common/b3Quaternion.h" #include "OpenGLWindow/b3gWindowInterface.h" #include "Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h" +#include "Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.h" + #include "../GpuDemoInternalData.h" #include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h" #include "OpenGLWindow/OpenGLInclude.h" @@ -103,7 +105,7 @@ __kernel void updateAabbSimple( __global float4* posOrnColors, const int numNode struct PairBenchInternalData { - b3GpuSapBroadphase* m_broadphaseGPU; + b3GpuBroadphaseInterface* m_broadphaseGPU; cl_kernel m_moveObjectsKernel; cl_kernel m_sineWaveKernel; @@ -176,6 +178,8 @@ void PairBench::initPhysics(const ConstructionInfo& ci) if (m_clData->m_clContext) { m_data->m_broadphaseGPU = new b3GpuSapBroadphase(m_clData->m_clContext,m_clData->m_clDevice,m_clData->m_clQueue); + //m_data->m_broadphaseGPU = new b3GpuGridBroadphase(m_clData->m_clContext,m_clData->m_clDevice,m_clData->m_clQueue); + cl_program pairBenchProg=0; int errNum=0; m_data->m_moveObjectsKernel = b3OpenCLUtils::compileCLKernelFromString(m_clData->m_clContext,m_clData->m_clDevice,s_pairBenchKernelString,"moveObjectsKernel",&errNum,pairBenchProg); @@ -446,7 +450,7 @@ void PairBench::clientMoveAndDisplay() } } - bool updateOnGpu=false; + bool updateOnGpu=false;//true; if (updateOnGpu) { @@ -463,7 +467,7 @@ void PairBench::clientMoveAndDisplay() B3_PROFILE("updateOnCpu"); if (!gPairBenchFileName) { - int allAabbs = m_data->m_broadphaseGPU->m_allAabbsCPU.size(); + int allAabbs = m_data->m_broadphaseGPU->getAllAabbsCPU().size(); b3AlignedObjectArray posOrnColorsCpu; @@ -476,10 +480,10 @@ void PairBench::clientMoveAndDisplay() { b3Vector3 position = posOrnColorsCpu[nodeId]; b3Vector3 halfExtents = b3MakeFloat4(1.01f,1.01f,1.01f,0.f); - m_data->m_broadphaseGPU->m_allAabbsCPU[nodeId].m_minVec = position-halfExtents; - m_data->m_broadphaseGPU->m_allAabbsCPU[nodeId].m_minIndices[3] = nodeId; - m_data->m_broadphaseGPU->m_allAabbsCPU[nodeId].m_maxVec = position+halfExtents; - m_data->m_broadphaseGPU->m_allAabbsCPU[nodeId].m_signedMaxIndices[3]= nodeId; + m_data->m_broadphaseGPU->getAllAabbsCPU()[nodeId].m_minVec = position-halfExtents; + m_data->m_broadphaseGPU->getAllAabbsCPU()[nodeId].m_minIndices[3] = nodeId; + m_data->m_broadphaseGPU->getAllAabbsCPU()[nodeId].m_maxVec = position+halfExtents; + m_data->m_broadphaseGPU->getAllAabbsCPU()[nodeId].m_signedMaxIndices[3]= nodeId; } } m_data->m_broadphaseGPU->writeAabbsToGpu(); @@ -506,7 +510,7 @@ void PairBench::clientMoveAndDisplay() if (m_data->m_gui) { - int allAabbs = m_data->m_broadphaseGPU->m_allAabbsCPU.size(); + int allAabbs = m_data->m_broadphaseGPU->getAllAabbsCPU().size(); int numOverlap = m_data->m_broadphaseGPU->getNumOverlap(); float time = dt/1000.f; diff --git a/Demos3/GpuDemos/main_opengl3core.cpp b/Demos3/GpuDemos/main_opengl3core.cpp index c6d545deb..54fc22541 100644 --- a/Demos3/GpuDemos/main_opengl3core.cpp +++ b/Demos3/GpuDemos/main_opengl3core.cpp @@ -100,7 +100,7 @@ GpuDemo::CreateFunc* allDemos[]= ConcaveScene::MyCreateFunc, - + GpuBoxPlaneScene::MyCreateFunc, GpuConstraintsDemo::MyCreateFunc, diff --git a/Demos3/GpuDemos/rigidbody/GpuRigidBodyDemo.cpp b/Demos3/GpuDemos/rigidbody/GpuRigidBodyDemo.cpp index 4d99cb339..26945c3d6 100644 --- a/Demos3/GpuDemos/rigidbody/GpuRigidBodyDemo.cpp +++ b/Demos3/GpuDemos/rigidbody/GpuRigidBodyDemo.cpp @@ -4,6 +4,8 @@ #include "Bullet3Common/b3Quaternion.h" #include "OpenGLWindow/b3gWindowInterface.h" #include "Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h" +#include "Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.h" + #include "../GpuDemoInternalData.h" #include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h" #include "OpenGLWindow/OpenGLInclude.h" @@ -119,6 +121,7 @@ void GpuRigidBodyDemo::initPhysics(const ConstructionInfo& ci) b3GpuNarrowPhase* np = new b3GpuNarrowPhase(m_clData->m_clContext,m_clData->m_clDevice,m_clData->m_clQueue,m_data->m_config); b3GpuSapBroadphase* bp = new b3GpuSapBroadphase(m_clData->m_clContext,m_clData->m_clDevice,m_clData->m_clQueue); + //b3GpuBroadphaseInterface* bp = new b3GpuGridBroadphase(m_clData->m_clContext,m_clData->m_clDevice,m_clData->m_clQueue); m_data->m_np = np; m_data->m_bp = bp; m_data->m_broadphaseDbvt = new b3DynamicBvhBroadphase(m_data->m_config.m_maxConvexBodies); @@ -419,6 +422,8 @@ bool GpuRigidBodyDemo::mouseButtonCallback(int button, int state, float x, float b3Quaternion orn(0,0,0,1); int fixedSphere = m_data->m_np->registerConvexHullShape(0,0,0,0);//>registerSphereShape(0.1); m_data->m_pickFixedBody = m_data->m_rigidBodyPipeline->registerPhysicsInstance(0,pos,orn,fixedSphere,0,false); + m_data->m_rigidBodyPipeline->writeAllInstancesToGpu(); + m_data->m_bp->writeAabbsToGpu(); if (m_data->m_pickGraphicsShapeIndex<0) { diff --git a/Demos3/GpuDemos/rigidbody/GpuRigidBodyDemoInternalData.h b/Demos3/GpuDemos/rigidbody/GpuRigidBodyDemoInternalData.h index f0c99e979..9147a5080 100644 --- a/Demos3/GpuDemos/rigidbody/GpuRigidBodyDemoInternalData.h +++ b/Demos3/GpuDemos/rigidbody/GpuRigidBodyDemoInternalData.h @@ -16,7 +16,7 @@ struct GpuRigidBodyDemoInternalData class b3GpuRigidBodyPipeline* m_rigidBodyPipeline; class b3GpuNarrowPhase* m_np; - class b3GpuSapBroadphase* m_bp; + class b3GpuBroadphaseInterface* m_bp; class b3DynamicBvhBroadphase* m_broadphaseDbvt; b3Vector3 m_pickPivotInA; diff --git a/Demos3/ImplicitCloth/premake4.lua b/Demos3/ImplicitCloth/premake4.lua index b767e95cc..8eea82ce3 100644 --- a/Demos3/ImplicitCloth/premake4.lua +++ b/Demos3/ImplicitCloth/premake4.lua @@ -8,6 +8,7 @@ includedirs { ".", + "../../src", "../../btgui" } @@ -39,8 +40,8 @@ "../../btgui/OpenGLTrueTypeFont/fontstash.h", "../../btgui/OpenGLTrueTypeFont/opengl_fontstashcallbacks.cpp", "../../btgui/OpenGLTrueTypeFont/opengl_fontstashcallbacks.h", - "../../btgui/Bullet3Common/**.cpp", - "../../btgui/Bullet3Common/**.h", + "../../src/Bullet3Common/**.cpp", + "../../src/Bullet3Common/**.h", "../../btgui/Timing/b3Clock.cpp", "../../btgui/Timing/b3Clock.h" diff --git a/Demos3/SimpleOpenGL3/premake4.lua b/Demos3/SimpleOpenGL3/premake4.lua index 4cc5a69d5..318e8caf5 100644 --- a/Demos3/SimpleOpenGL3/premake4.lua +++ b/Demos3/SimpleOpenGL3/premake4.lua @@ -8,6 +8,7 @@ includedirs { ".", + "../../src", "../../btgui" } @@ -15,10 +16,12 @@ initGlew() --links{"gwen"} - + files { "**.cpp", "**.h", + "../../src/Bullet3Common/**.cpp", + "../../src/Bullet3Common/**.h", "../../btgui/OpenGLWindow/SimpleOpenGL3App.cpp", "../../btgui/OpenGLWindow/SimpleOpenGL3App.h", "../../btgui/OpenGLWindow/TwFonts.cpp", diff --git a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuBroadphaseInterface.h b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuBroadphaseInterface.h new file mode 100644 index 000000000..7e5ba32ee --- /dev/null +++ b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuBroadphaseInterface.h @@ -0,0 +1,39 @@ + +#ifndef B3_GPU_BROADPHASE_INTERFACE_H +#define B3_GPU_BROADPHASE_INTERFACE_H + +#include "Bullet3OpenCL/Initialize/b3OpenCLInclude.h" +#include "Bullet3Common/b3Vector3.h" +#include "b3SapAabb.h" +#include "Bullet3Common/shared/b3Int2.h" +#include "Bullet3Common/shared/b3Int4.h" +#include "Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h" + +class b3GpuBroadphaseInterface +{ +public: + + + virtual ~b3GpuBroadphaseInterface() + { + } + + virtual void createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask)=0; + virtual void createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask)=0; + + virtual void calculateOverlappingPairs(int maxPairs)=0; + virtual void calculateOverlappingPairsHost(int maxPairs)=0; + + //call writeAabbsToGpu after done making all changes (createProxy etc) + virtual void writeAabbsToGpu()=0; + + virtual cl_mem getAabbBufferWS()=0; + virtual int getNumOverlap()=0; + virtual cl_mem getOverlappingPairBuffer()=0; + + virtual b3OpenCLArray& getAllAabbsGPU()=0; + virtual b3AlignedObjectArray& getAllAabbsCPU()=0; + +}; + +#endif //B3_GPU_BROADPHASE_INTERFACE_H diff --git a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.cpp b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.cpp new file mode 100644 index 000000000..7e022d3d8 --- /dev/null +++ b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.cpp @@ -0,0 +1,98 @@ + +#include "b3GpuGridBroadphase.h" +#include "Bullet3Geometry/b3AabbUtil.h" + +b3GpuGridBroadphase::b3GpuGridBroadphase(cl_context ctx,cl_device_id device, cl_command_queue q ) +:m_context(ctx), +m_device(device), +m_queue(q), +m_allAabbsGPU(ctx,q), +m_gpuPairs(ctx,q) +{ +} +b3GpuGridBroadphase::~b3GpuGridBroadphase() +{ +} + + + +void b3GpuGridBroadphase::createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask) +{ + b3SapAabb aabb; + aabb.m_minVec = aabbMin; + aabb.m_maxVec = aabbMax; + aabb.m_minIndices[3] = userPtr; + aabb.m_signedMaxIndices[3] = userPtr; + m_allAabbsCPU.push_back(aabb); +} +void b3GpuGridBroadphase::createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask) +{ + createProxy(aabbMin,aabbMax,userPtr,collisionFilterGroup,collisionFilterMask); + +} + +void b3GpuGridBroadphase::calculateOverlappingPairs(int maxPairs) +{ + calculateOverlappingPairsHost(maxPairs); +} +void b3GpuGridBroadphase::calculateOverlappingPairsHost(int maxPairs) +{ + m_hostPairs.resize(0); + + for (int i=0;im_allAabbsGPU.getBufferCL(); +} +int b3GpuGridBroadphase::getNumOverlap() +{ + return m_gpuPairs.size(); +} +cl_mem b3GpuGridBroadphase::getOverlappingPairBuffer() +{ + return m_gpuPairs.getBufferCL(); +} + +b3OpenCLArray& b3GpuGridBroadphase::getAllAabbsGPU() +{ + return m_allAabbsGPU; +} + +b3AlignedObjectArray& b3GpuGridBroadphase::getAllAabbsCPU() +{ + return m_allAabbsCPU; +} \ No newline at end of file diff --git a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.h b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.h new file mode 100644 index 000000000..8bd21cbcb --- /dev/null +++ b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.h @@ -0,0 +1,44 @@ +#ifndef B3_GPU_GRID_BROADPHASE_H +#define B3_GPU_GRID_BROADPHASE_H + +#include "b3GpuBroadphaseInterface.h" + +class b3GpuGridBroadphase : public b3GpuBroadphaseInterface +{ +protected: + cl_context m_context; + cl_device_id m_device; + cl_command_queue m_queue; + + b3OpenCLArray m_allAabbsGPU; + b3AlignedObjectArray m_allAabbsCPU; + + b3AlignedObjectArray m_hostPairs; + b3OpenCLArray m_gpuPairs; + +public: + + b3GpuGridBroadphase(cl_context ctx,cl_device_id device, cl_command_queue q ); + virtual ~b3GpuGridBroadphase(); + + + + virtual void createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask); + virtual void createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask); + + virtual void calculateOverlappingPairs(int maxPairs); + virtual void calculateOverlappingPairsHost(int maxPairs); + + //call writeAabbsToGpu after done making all changes (createProxy etc) + virtual void writeAabbsToGpu(); + + virtual cl_mem getAabbBufferWS(); + virtual int getNumOverlap(); + virtual cl_mem getOverlappingPairBuffer(); + + virtual b3OpenCLArray& getAllAabbsGPU(); + virtual b3AlignedObjectArray& getAllAabbsCPU(); + +}; + +#endif //B3_GPU_GRID_BROADPHASE_H \ No newline at end of file diff --git a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h index 7d113f23c..afcc93a83 100644 --- a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h +++ b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h @@ -9,8 +9,9 @@ class b3Vector3; #include "b3SapAabb.h" #include "Bullet3Common/shared/b3Int2.h" +#include "b3GpuBroadphaseInterface.h" -class b3GpuSapBroadphase +class b3GpuSapBroadphase : public b3GpuBroadphaseInterface { cl_context m_context; @@ -58,6 +59,15 @@ class b3GpuSapBroadphase b3OpenCLArray m_allAabbsGPU; b3AlignedObjectArray m_allAabbsCPU; + virtual b3OpenCLArray& getAllAabbsGPU() + { + return m_allAabbsGPU; + } + virtual b3AlignedObjectArray& getAllAabbsCPU() + { + return m_allAabbsCPU; + } + b3OpenCLArray m_sum; b3OpenCLArray m_sum2; b3OpenCLArray m_dst; @@ -79,23 +89,23 @@ class b3GpuSapBroadphase b3GpuSapBroadphase(cl_context ctx,cl_device_id device, cl_command_queue q ); virtual ~b3GpuSapBroadphase(); - void calculateOverlappingPairs(int maxPairs); - void calculateOverlappingPairsHost(int maxPairs); + virtual void calculateOverlappingPairs(int maxPairs); + virtual void calculateOverlappingPairsHost(int maxPairs); void reset(); void init3dSap(); - void calculateOverlappingPairsHostIncremental3Sap(); + virtual void calculateOverlappingPairsHostIncremental3Sap(); - void createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask); - void createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask); + virtual void createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask); + virtual void createLargeProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask); //call writeAabbsToGpu after done making all changes (createProxy etc) - void writeAabbsToGpu(); + virtual void writeAabbsToGpu(); - cl_mem getAabbBufferWS(); - int getNumOverlap(); - cl_mem getOverlappingPairBuffer(); + virtual cl_mem getAabbBufferWS(); + virtual int getNumOverlap(); + virtual cl_mem getOverlappingPairBuffer(); }; #endif //B3_GPU_SAP_BROADPHASE_H \ No newline at end of file diff --git a/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.cpp b/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.cpp index dbe81d500..fcd95ae07 100644 --- a/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.cpp +++ b/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.cpp @@ -22,7 +22,7 @@ subject to the following restrictions: #include "b3GpuNarrowPhase.h" #include "Bullet3Geometry/b3AabbUtil.h" #include "Bullet3OpenCL/BroadphaseCollision/b3SapAabb.h" -#include "Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h" +#include "Bullet3OpenCL/BroadphaseCollision/b3GpuBroadphaseInterface.h" #include "Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h" #include "Bullet3Dynamics/ConstraintSolver/b3PgsJacobiSolver.h" #include "Bullet3Collision/NarrowPhaseCollision/shared/b3UpdateAabbs.h" @@ -43,7 +43,7 @@ subject to the following restrictions: bool integrateOnCpu = true; #else - bool useDbvt = false;//true; + bool useDbvt = false; bool useBullet2CpuSolver = true; bool dumpContactStats = false; bool calcWorldSpaceAabbOnCpu = false;//true; @@ -70,7 +70,7 @@ subject to the following restrictions: #include "Bullet3Dynamics/shared/b3IntegrateTransforms.h" #include "Bullet3OpenCL/RigidBody/b3GpuNarrowPhaseInternalData.h" -b3GpuRigidBodyPipeline::b3GpuRigidBodyPipeline(cl_context ctx,cl_device_id device, cl_command_queue q,class b3GpuNarrowPhase* narrowphase, class b3GpuSapBroadphase* broadphaseSap , struct b3DynamicBvhBroadphase* broadphaseDbvt, const b3Config& config) +b3GpuRigidBodyPipeline::b3GpuRigidBodyPipeline(cl_context ctx,cl_device_id device, cl_command_queue q,class b3GpuNarrowPhase* narrowphase, class b3GpuBroadphaseInterface* broadphaseSap , struct b3DynamicBvhBroadphase* broadphaseDbvt, const b3Config& config) { m_data = new b3GpuRigidBodyPipelineInternalData; m_data->m_constraintUid=0; @@ -497,13 +497,14 @@ void b3GpuRigidBodyPipeline::setupGpuAabbsFull() m_data->m_allAabbsGPU->copyFromHost(m_data->m_allAabbsCPU); } else { - m_data->m_broadphaseSap->m_allAabbsCPU.resize(numBodies); + m_data->m_broadphaseSap->getAllAabbsCPU().resize(numBodies); m_data->m_narrowphase->readbackAllBodiesToCpu(); for (int i=0;im_narrowphase->getBodiesCpu(), m_data->m_narrowphase->getCollidablesCpu(), m_data->m_narrowphase->getLocalSpaceAabbsCpu(),&m_data->m_broadphaseSap->m_allAabbsCPU[0]); + b3ComputeWorldAabb( i, m_data->m_narrowphase->getBodiesCpu(), m_data->m_narrowphase->getCollidablesCpu(), m_data->m_narrowphase->getLocalSpaceAabbsCpu(),&m_data->m_broadphaseSap->getAllAabbsCPU()[0]); } - m_data->m_broadphaseSap->m_allAabbsGPU.copyFromHost(m_data->m_broadphaseSap->m_allAabbsCPU); + m_data->m_broadphaseSap->getAllAabbsGPU().copyFromHost(m_data->m_broadphaseSap->getAllAabbsCPU()); + //m_data->m_broadphaseSap->writeAabbsToGpu(); } } } else @@ -631,10 +632,10 @@ int b3GpuRigidBodyPipeline::registerPhysicsInstance(float mass, const float* po { if (mass) { - m_data->m_broadphaseSap->createProxy(aabbMin,aabbMax,userIndex,1,1);//m_dispatcher); + m_data->m_broadphaseSap->createProxy(aabbMin,aabbMax,bodyIndex,1,1);//m_dispatcher); } else { - m_data->m_broadphaseSap->createLargeProxy(aabbMin,aabbMax,userIndex,1,1);//m_dispatcher); + m_data->m_broadphaseSap->createLargeProxy(aabbMin,aabbMax,bodyIndex,1,1);//m_dispatcher); } } } diff --git a/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.h b/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.h index ffbd8b3f9..b4eac6841 100644 --- a/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.h +++ b/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.h @@ -32,7 +32,7 @@ protected: public: - b3GpuRigidBodyPipeline(cl_context ctx,cl_device_id device, cl_command_queue q , class b3GpuNarrowPhase* narrowphase, class b3GpuSapBroadphase* broadphaseSap, struct b3DynamicBvhBroadphase* broadphaseDbvt, const b3Config& config); + b3GpuRigidBodyPipeline(cl_context ctx,cl_device_id device, cl_command_queue q , class b3GpuNarrowPhase* narrowphase, class b3GpuBroadphaseInterface* broadphaseSap, struct b3DynamicBvhBroadphase* broadphaseDbvt, const b3Config& config); virtual ~b3GpuRigidBodyPipeline(); void stepSimulation(float deltaTime); diff --git a/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipelineInternalData.h b/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipelineInternalData.h index e798b1cf6..533c94ce9 100644 --- a/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipelineInternalData.h +++ b/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipelineInternalData.h @@ -50,7 +50,7 @@ struct b3GpuRigidBodyPipelineInternalData class b3GpuJacobiSolver* m_solver3; class b3GpuRaycast* m_raycaster; - class b3GpuSapBroadphase* m_broadphaseSap; + class b3GpuBroadphaseInterface* m_broadphaseSap; struct b3DynamicBvhBroadphase* m_broadphaseDbvt; b3OpenCLArray* m_allAabbsGPU;