From 9e623f6b0f4dc7c85f2a8cc0f07cf5f6e9ad2ce1 Mon Sep 17 00:00:00 2001 From: Erwin Coumans Date: Wed, 13 Mar 2013 15:03:08 -0700 Subject: [PATCH] add broadphase demo (empty skeleton) --- demo/gpudemo/GpuDemo.cpp | 79 +++++++++++++++++++ demo/gpudemo/GpuDemo.h | 64 +++++++++++++++ demo/gpudemo/GpuDemoInternalData.h | 22 ++++++ demo/gpudemo/ParticleDemo.cpp | 107 ++++++++------------------ demo/gpudemo/ParticleDemo.h | 46 ++--------- demo/gpudemo/broadphase/PairBench.cpp | 23 ++++++ demo/gpudemo/broadphase/PairBench.h | 34 ++++++++ demo/gpudemo/main_opengl3core.cpp | 9 ++- demo/gpudemo/premake4.lua | 7 +- 9 files changed, 267 insertions(+), 124 deletions(-) create mode 100644 demo/gpudemo/GpuDemo.cpp create mode 100644 demo/gpudemo/GpuDemo.h create mode 100644 demo/gpudemo/GpuDemoInternalData.h create mode 100644 demo/gpudemo/broadphase/PairBench.cpp create mode 100644 demo/gpudemo/broadphase/PairBench.h diff --git a/demo/gpudemo/GpuDemo.cpp b/demo/gpudemo/GpuDemo.cpp new file mode 100644 index 000000000..c4c845b66 --- /dev/null +++ b/demo/gpudemo/GpuDemo.cpp @@ -0,0 +1,79 @@ +#include "GpuDemo.h" +#include "GpuDemoInternalData.h" +#include "BulletCommon/btScalar.h" +#include "basic_initialize/btOpenCLUtils.h" + + +GpuDemo::GpuDemo() +:m_clData(0) +{ + m_clData = new GpuDemoInternalData(); +} + +GpuDemo::~GpuDemo() +{ + if (m_clData) + { + btAssert(m_clData->m_clInitialized==false); + + delete m_clData; + } +} + +void GpuDemo::exitCL() +{ + if (m_clData->m_clInitialized) + { + clReleaseCommandQueue(m_clData->m_clQueue); + clReleaseContext(m_clData->m_clContext); + m_clData->m_clInitialized = false; + } + +} + +void GpuDemo::initCL(int preferredDeviceIndex, int preferredPlatformIndex) +{ + void* glCtx=0; + void* glDC = 0; + + + + int ciErrNum = 0; + //#ifdef CL_PLATFORM_INTEL + // cl_device_type deviceType = CL_DEVICE_TYPE_ALL; + //#else + cl_device_type deviceType = CL_DEVICE_TYPE_GPU; + //#endif + + + + // if (useInterop) + // { + // m_data->m_clContext = btOpenCLUtils::createContextFromType(deviceType, &ciErrNum, glCtx, glDC); + // } else + { + m_clData->m_clContext = btOpenCLUtils::createContextFromType(deviceType, &ciErrNum, 0,0,preferredDeviceIndex, preferredPlatformIndex); + } + + + oclCHECKERROR(ciErrNum, CL_SUCCESS); + + int numDev = btOpenCLUtils::getNumDevices(m_clData->m_clContext); + + if (numDev>0) + { + m_clData->m_clDevice= btOpenCLUtils::getDevice(m_clData->m_clContext,0); + m_clData->m_clQueue = clCreateCommandQueue(m_clData->m_clContext, m_clData->m_clDevice, 0, &ciErrNum); + oclCHECKERROR(ciErrNum, CL_SUCCESS); + + btOpenCLUtils::printDeviceInfo(m_clData->m_clDevice); + btOpenCLDeviceInfo info; + btOpenCLUtils::getDeviceInfo(m_clData->m_clDevice,&info); + m_clData->m_clDeviceName = info.m_deviceName; + m_clData->m_clInitialized = true; + + } + +} + + diff --git a/demo/gpudemo/GpuDemo.h b/demo/gpudemo/GpuDemo.h new file mode 100644 index 000000000..33cff011c --- /dev/null +++ b/demo/gpudemo/GpuDemo.h @@ -0,0 +1,64 @@ +#ifndef GPU_DEMO_H +#define GPU_DEMO_H +class GLInstancingRenderer; + +class GpuDemo +{ +protected: + + struct GpuDemoInternalData* m_clData; + + + virtual void initCL(int preferredDeviceIndex, int preferredPlatformIndex); + virtual void exitCL(); +public: + + typedef class GpuDemo* (CreateFunc)(); + + struct ConstructionInfo + { + bool useOpenCL; + int preferredOpenCLPlatformIndex; + int preferredOpenCLDeviceIndex; + int arraySizeX; + int arraySizeY; + int arraySizeZ; + bool m_useConcaveMesh; + float gapX; + float gapY; + float gapZ; + GLInstancingRenderer* m_instancingRenderer; + ConstructionInfo() + :useOpenCL(false),//true), + preferredOpenCLPlatformIndex(-1), + preferredOpenCLDeviceIndex(-1), + arraySizeX(10), + arraySizeY(10 ), + arraySizeZ(10), + m_useConcaveMesh(false), + gapX(4.3), + gapY(4.0), + gapZ(4.3), + m_instancingRenderer(0) + { + } + }; + + GpuDemo(); + virtual ~GpuDemo(); + + virtual const char* getName()=0; + + virtual void initPhysics(const ConstructionInfo& ci)=0; + + virtual void exitPhysics()=0; + + virtual void renderScene()=0; + + virtual void clientMoveAndDisplay()=0; + + +}; + +#endif + diff --git a/demo/gpudemo/GpuDemoInternalData.h b/demo/gpudemo/GpuDemoInternalData.h new file mode 100644 index 000000000..d35aa9edf --- /dev/null +++ b/demo/gpudemo/GpuDemoInternalData.h @@ -0,0 +1,22 @@ +#ifndef GPU_DEMO_INTERNAL_DATA_H +#define GPU_DEMO_INTERNAL_DATA_H + +#include "basic_initialize/btOpenCLInclude.h" + +struct GpuDemoInternalData +{ + cl_context m_clContext; + cl_device_id m_clDevice; + cl_command_queue m_clQueue; + bool m_clInitialized; + char* m_clDeviceName; + + GpuDemoInternalData() + :m_clInitialized(false), + m_clDeviceName(0) + { + + } +}; + +#endif diff --git a/demo/gpudemo/ParticleDemo.cpp b/demo/gpudemo/ParticleDemo.cpp index bfab839ec..0f2cad18d 100644 --- a/demo/gpudemo/ParticleDemo.cpp +++ b/demo/gpudemo/ParticleDemo.cpp @@ -16,7 +16,7 @@ static char* particleKernelsString = //#include "../../opencl/primitives/AdlPrimitives/Math/Math.h" //#include "../../opencl/broadphase_benchmark/btGridBroadphaseCL.h" #include "gpu_broadphase/host/btGpuSapBroadphase.h" - +#include "GpuDemoInternalData.h" #include "BulletCommon/btQuickprof.h" @@ -84,9 +84,7 @@ ATTRIBUTE_ALIGNED16(struct) btSimParams struct ParticleInternalData { - cl_context m_clContext; - cl_device_id m_clDevice; - cl_command_queue m_clQueue; + cl_kernel m_updatePositionsKernel; cl_kernel m_updatePositionsKernel2; @@ -105,10 +103,10 @@ struct ParticleInternalData btAlignedObjectArray m_simParamCPU; btOpenCLArray* m_simParamGPU; - bool m_clInitialized; + ParticleInternalData() - :m_clInitialized(false), + : m_clPositionBuffer(0), m_velocitiesGPU(0), m_simParamGPU(0), @@ -120,7 +118,6 @@ struct ParticleInternalData m_simParamCPU.resize(1); } - char* m_clDeviceName; }; @@ -141,62 +138,20 @@ ParticleDemo::~ParticleDemo() void ParticleDemo::exitCL() { - if (m_data->m_clInitialized) + if (m_clData->m_clInitialized) { - m_data->m_clInitialized = false; - clReleaseCommandQueue(m_data->m_clQueue); clReleaseKernel(m_data->m_updatePositionsKernel); clReleaseKernel(m_data->m_updatePositionsKernel2); clReleaseKernel(m_data->m_updateAabbsKernel); clReleaseKernel(m_data->m_collideParticlesKernel); - - clReleaseContext(m_data->m_clContext); } + + GpuDemo::exitCL(); } void ParticleDemo::initCL(int preferredDeviceIndex, int preferredPlatformIndex) { - void* glCtx=0; - void* glDC = 0; - - - - int ciErrNum = 0; -//#ifdef CL_PLATFORM_INTEL -// cl_device_type deviceType = CL_DEVICE_TYPE_ALL; -//#else - cl_device_type deviceType = CL_DEVICE_TYPE_GPU; -//#endif - - - -// if (useInterop) -// { -// m_data->m_clContext = btOpenCLUtils::createContextFromType(deviceType, &ciErrNum, glCtx, glDC); -// } else - { - m_data->m_clContext = btOpenCLUtils::createContextFromType(deviceType, &ciErrNum, 0,0,preferredDeviceIndex, preferredPlatformIndex); - } - - - oclCHECKERROR(ciErrNum, CL_SUCCESS); - - int numDev = btOpenCLUtils::getNumDevices(m_data->m_clContext); - - if (numDev>0) - { - m_data->m_clDevice= btOpenCLUtils::getDevice(m_data->m_clContext,0); - m_data->m_clQueue = clCreateCommandQueue(m_data->m_clContext, m_data->m_clDevice, 0, &ciErrNum); - oclCHECKERROR(ciErrNum, CL_SUCCESS); - - btOpenCLUtils::printDeviceInfo(m_data->m_clDevice); - btOpenCLDeviceInfo info; - btOpenCLUtils::getDeviceInfo(m_data->m_clDevice,&info); - m_data->m_clDeviceName = info.m_deviceName; - m_data->m_clInitialized = true; - - } - + GpuDemo::initCL(preferredDeviceIndex,preferredPlatformIndex); } @@ -213,14 +168,14 @@ void ParticleDemo::setupScene(const ConstructionInfo& ci) int maxPairsSmallProxy = 32; float radius = 3.f*m_data->m_simParamCPU[0].m_particleRad; - m_data->m_broadphaseGPU = new btGpuSapBroadphase(m_data->m_clContext ,m_data->m_clDevice,m_data->m_clQueue);//overlappingPairCache,btVector3(4.f, 4.f, 4.f), 128, 128, 128,maxObjects, maxObjects, maxPairsSmallProxy, 100.f, 128, + m_data->m_broadphaseGPU = new btGpuSapBroadphase(m_clData->m_clContext ,m_clData->m_clDevice,m_clData->m_clQueue);//overlappingPairCache,btVector3(4.f, 4.f, 4.f), 128, 128, 128,maxObjects, maxObjects, maxPairsSmallProxy, 100.f, 128, /*m_data->m_broadphaseGPU = new btGridBroadphaseCl(overlappingPairCache,btVector3(radius,radius,radius), 128, 128, 128, maxObjects, maxObjects, maxPairsSmallProxy, 100.f, 128, - m_data->m_clContext ,m_data->m_clDevice,m_data->m_clQueue); + m_clData->m_clContext ,m_clData->m_clDevice,m_clData->m_clQueue); */ - m_data->m_velocitiesGPU = new btOpenCLArray(m_data->m_clContext,m_data->m_clQueue,numParticles); + m_data->m_velocitiesGPU = new btOpenCLArray(m_clData->m_clContext,m_clData->m_clQueue,numParticles); m_data->m_velocitiesCPU.resize(numParticles); for (int i=0;im_velocitiesGPU->copyFromHost(m_data->m_velocitiesCPU); - m_data->m_simParamGPU = new btOpenCLArray(m_data->m_clContext,m_data->m_clQueue,1,false); + m_data->m_simParamGPU = new btOpenCLArray(m_clData->m_clContext,m_clData->m_clQueue,1,false); m_data->m_simParamGPU->copyFromHost(m_data->m_simParamCPU); cl_int pErrNum; - cl_program prog = btOpenCLUtils::compileCLProgramFromString(m_data->m_clContext,m_data->m_clDevice,particleKernelsString,0,"",INTEROPKERNEL_SRC_PATH); - m_data->m_updatePositionsKernel = btOpenCLUtils::compileCLKernelFromString(m_data->m_clContext, m_data->m_clDevice,particleKernelsString, "updatePositionsKernel" ,&pErrNum,prog); + cl_program prog = btOpenCLUtils::compileCLProgramFromString(m_clData->m_clContext,m_clData->m_clDevice,particleKernelsString,0,"",INTEROPKERNEL_SRC_PATH); + m_data->m_updatePositionsKernel = btOpenCLUtils::compileCLKernelFromString(m_clData->m_clContext, m_clData->m_clDevice,particleKernelsString, "updatePositionsKernel" ,&pErrNum,prog); oclCHECKERROR(pErrNum, CL_SUCCESS); - m_data->m_updatePositionsKernel2 = btOpenCLUtils::compileCLKernelFromString(m_data->m_clContext, m_data->m_clDevice,particleKernelsString, "integrateMotionKernel" ,&pErrNum,prog); + m_data->m_updatePositionsKernel2 = btOpenCLUtils::compileCLKernelFromString(m_clData->m_clContext, m_clData->m_clDevice,particleKernelsString, "integrateMotionKernel" ,&pErrNum,prog); oclCHECKERROR(pErrNum, CL_SUCCESS); - m_data->m_updateAabbsKernel= btOpenCLUtils::compileCLKernelFromString(m_data->m_clContext, m_data->m_clDevice,particleKernelsString, "updateAabbsKernel" ,&pErrNum,prog); + m_data->m_updateAabbsKernel= btOpenCLUtils::compileCLKernelFromString(m_clData->m_clContext, m_clData->m_clDevice,particleKernelsString, "updateAabbsKernel" ,&pErrNum,prog); oclCHECKERROR(pErrNum, CL_SUCCESS); - m_data->m_collideParticlesKernel = btOpenCLUtils::compileCLKernelFromString(m_data->m_clContext, m_data->m_clDevice,particleKernelsString, "collideParticlesKernel" ,&pErrNum,prog); + m_data->m_collideParticlesKernel = btOpenCLUtils::compileCLKernelFromString(m_clData->m_clContext, m_clData->m_clDevice,particleKernelsString, "collideParticlesKernel" ,&pErrNum,prog); oclCHECKERROR(pErrNum, CL_SUCCESS); m_instancingRenderer = ci.m_instancingRenderer; @@ -368,15 +323,15 @@ void ParticleDemo::clientMoveAndDisplay() cl_int ciErrNum; if (!m_data->m_clPositionBuffer) { - m_data->m_clPositionBuffer = clCreateBuffer(m_data->m_clContext, CL_MEM_READ_WRITE, + m_data->m_clPositionBuffer = clCreateBuffer(m_clData->m_clContext, CL_MEM_READ_WRITE, posArraySize, 0, &ciErrNum); - clFinish(m_data->m_clQueue); + clFinish(m_clData->m_clQueue); oclCHECKERROR(ciErrNum, CL_SUCCESS); - ciErrNum = clEnqueueWriteBuffer ( m_data->m_clQueue,m_data->m_clPositionBuffer, + ciErrNum = clEnqueueWriteBuffer ( m_clData->m_clQueue,m_data->m_clPositionBuffer, blocking,0,posArraySize,hostPtr,0,0,0 ); - clFinish(m_data->m_clQueue); + clFinish(m_clData->m_clQueue); } @@ -391,13 +346,13 @@ void ParticleDemo::clientMoveAndDisplay() btBufferInfoCL( m_data->m_clPositionBuffer) }; - btLauncherCL launcher(m_data->m_clQueue, m_data->m_updatePositionsKernel ); + btLauncherCL launcher(m_clData->m_clQueue, m_data->m_updatePositionsKernel ); launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(btBufferInfoCL) ); launcher.setConst( numParticles); launcher.launch1D( numParticles); - clFinish(m_data->m_clQueue); + clFinish(m_clData->m_clQueue); } @@ -410,7 +365,7 @@ void ParticleDemo::clientMoveAndDisplay() btBufferInfoCL( m_data->m_simParamGPU->getBufferCL(),true) }; - btLauncherCL launcher(m_data->m_clQueue, m_data->m_updatePositionsKernel2 ); + btLauncherCL launcher(m_clData->m_clQueue, m_data->m_updatePositionsKernel2 ); launcher.setConst( numParticles); launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(btBufferInfoCL) ); @@ -418,7 +373,7 @@ void ParticleDemo::clientMoveAndDisplay() launcher.setConst( timeStep); launcher.launch1D( numParticles); - clFinish(m_data->m_clQueue); + clFinish(m_clData->m_clQueue); } @@ -428,13 +383,13 @@ void ParticleDemo::clientMoveAndDisplay() btBufferInfoCL( m_data->m_broadphaseGPU->getAabbBuffer()), }; - btLauncherCL launcher(m_data->m_clQueue, m_data->m_updateAabbsKernel ); + btLauncherCL launcher(m_clData->m_clQueue, m_data->m_updateAabbsKernel ); launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(btBufferInfoCL) ); launcher.setConst( m_data->m_simParamCPU[0].m_particleRad); launcher.setConst( numParticles); launcher.launch1D( numParticles); - clFinish(m_data->m_clQueue); + clFinish(m_clData->m_clQueue); } //broadphase @@ -455,11 +410,11 @@ void ParticleDemo::clientMoveAndDisplay() btBufferInfoCL( m_data->m_broadphaseGPU->getOverlappingPairBuffer(),true), }; - btLauncherCL launcher(m_data->m_clQueue, m_data->m_collideParticlesKernel); + btLauncherCL launcher(m_clData->m_clQueue, m_data->m_collideParticlesKernel); launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(btBufferInfoCL) ); launcher.setConst( numPairsGPU); launcher.launch1D( numPairsGPU); - clFinish(m_data->m_clQueue); + clFinish(m_clData->m_clQueue); //__kernel void collideParticlesKernel( __global float4* pPos, __global float4* pVel, __global int2* pairs, const int numPairs) } @@ -467,7 +422,7 @@ void ParticleDemo::clientMoveAndDisplay() if (1) { - ciErrNum = clEnqueueReadBuffer ( m_data->m_clQueue, + ciErrNum = clEnqueueReadBuffer ( m_clData->m_clQueue, m_data->m_clPositionBuffer, blocking, 0, @@ -475,7 +430,7 @@ void ParticleDemo::clientMoveAndDisplay() hostPtr,0,0,0); //clReleaseMemObject(clBuffer); - clFinish(m_data->m_clQueue); + clFinish(m_clData->m_clQueue); } diff --git a/demo/gpudemo/ParticleDemo.h b/demo/gpudemo/ParticleDemo.h index a2e836032..228101c79 100644 --- a/demo/gpudemo/ParticleDemo.h +++ b/demo/gpudemo/ParticleDemo.h @@ -1,55 +1,21 @@ #ifndef PARTICLE_DEMO_H #define PARTICLE_DEMO_H -//#include "GpuDemo.h" +#include "GpuDemo.h" struct GLInstancingRenderer; -class ParticleDemo; -class ParticleDemo //: public GpuDemo +class ParticleDemo : public GpuDemo { public: - typedef class ParticleDemo* (CreateFunc)(); - - - - - struct ConstructionInfo - { - bool useOpenCL; - int preferredOpenCLPlatformIndex; - int preferredOpenCLDeviceIndex; - int arraySizeX; - int arraySizeY; - int arraySizeZ; - bool m_useConcaveMesh; - float gapX; - float gapY; - float gapZ; - GLInstancingRenderer* m_instancingRenderer; - ConstructionInfo() - :useOpenCL(false),//true), - preferredOpenCLPlatformIndex(-1), - preferredOpenCLDeviceIndex(-1), - arraySizeX(10), - arraySizeY(10 ), - arraySizeZ(10), - m_useConcaveMesh(false), - gapX(4.3), - gapY(4.0), - gapZ(4.3), - m_instancingRenderer(0) - { - } - }; protected: struct ParticleInternalData* m_data; GLInstancingRenderer* m_instancingRenderer; - void initCL(int preferredDeviceIndex, int preferredPlatformIndex); - void exitCL(); + virtual void initCL(int preferredDeviceIndex, int preferredPlatformIndex); + virtual void exitCL(); public: @@ -67,9 +33,9 @@ public: { return "ParticleDemo"; } - static ParticleDemo* MyCreateFunc() + static GpuDemo* MyCreateFunc() { - ParticleDemo* demo = new ParticleDemo; + GpuDemo* demo = new ParticleDemo; return demo; } diff --git a/demo/gpudemo/broadphase/PairBench.cpp b/demo/gpudemo/broadphase/PairBench.cpp new file mode 100644 index 000000000..27cb3c1df --- /dev/null +++ b/demo/gpudemo/broadphase/PairBench.cpp @@ -0,0 +1,23 @@ +#include "PairBench.h" +#include "BulletCommon/btQuickprof.h" + +void PairBench::initPhysics(const ConstructionInfo& ci) +{ + CProfileManager::CleanupMemory(); +} + +void PairBench::exitPhysics() +{ + +} + + +void PairBench::renderScene() +{ + +} + +void PairBench::clientMoveAndDisplay() +{ + +} diff --git a/demo/gpudemo/broadphase/PairBench.h b/demo/gpudemo/broadphase/PairBench.h new file mode 100644 index 000000000..bf688a88e --- /dev/null +++ b/demo/gpudemo/broadphase/PairBench.h @@ -0,0 +1,34 @@ +#ifndef PAIR_BENCH_H +#define PAIR_BENCH_H + +#include "../GpuDemo.h" + +class PairBench : public GpuDemo +{ +public: + + virtual void initPhysics(const ConstructionInfo& ci); + + virtual void exitPhysics(); + + virtual const char* getName() + { + return "PairBench"; + } + static GpuDemo* MyCreateFunc() + { + GpuDemo* demo = new PairBench; + return demo; + } + + + + virtual void renderScene(); + + virtual void clientMoveAndDisplay(); + + +}; + +#endif + diff --git a/demo/gpudemo/main_opengl3core.cpp b/demo/gpudemo/main_opengl3core.cpp index af66698a5..8bf24a8d9 100644 --- a/demo/gpudemo/main_opengl3core.cpp +++ b/demo/gpudemo/main_opengl3core.cpp @@ -24,6 +24,8 @@ #include "OpenGLTrueTypeFont/opengl_fontstashcallbacks.h" #include "gwenUserInterface.h" #include "ParticleDemo.h" +#include "broadphase/PairBench.h" + //#include "BroadphaseBenchmark.h" int g_OpenGLWidth=1024; @@ -56,11 +58,12 @@ enum btAlignedObjectArray demoNames; int selectedDemo = 0; -ParticleDemo::CreateFunc* allDemos[]= +GpuDemo::CreateFunc* allDemos[]= { //BroadphaseBenchmark::CreateFunc, //GpuBoxDemo::CreateFunc, ParticleDemo::MyCreateFunc, + PairBench::MyCreateFunc, //SpheresDemo::CreateFunc, //GpuCompoundDemo::CreateFunc, //EmptyDemo::CreateFunc, @@ -434,7 +437,7 @@ int main(int argc, char* argv[]) demoNames.clear(); for (int i=0;igetName()); delete demo; } @@ -526,7 +529,7 @@ int main(int argc, char* argv[]) { - ParticleDemo* demo = allDemos[selectedDemo](); + GpuDemo* demo = allDemos[selectedDemo](); // demo->myinit(); bool useGpu = false; diff --git a/demo/gpudemo/premake4.lua b/demo/gpudemo/premake4.lua index 6997bbca4..38b9ee4ed 100644 --- a/demo/gpudemo/premake4.lua +++ b/demo/gpudemo/premake4.lua @@ -29,11 +29,8 @@ function createProject(vendor) } files { - "main_opengl3core.cpp", - "gwenUserInterface.cpp", - "gwenUserInterface.h", - "ParticleDemo.cpp", - "ParticleDemo.h", + "**.cpp", + "**.h", "../../src/BulletCommon/btAlignedAllocator.cpp", "../../src/BulletCommon/btAlignedAllocator.h", "../../src/BulletCommon/btQuickprof.cpp",