From 39dbb51f683683c75abdf3ef5dfab5e86aeedd0a Mon Sep 17 00:00:00 2001 From: erwin coumans Date: Thu, 14 Mar 2013 14:35:19 -0700 Subject: [PATCH] improved PairBench. --- demo/gpudemo/GpuDemo.cpp | 5 +- demo/gpudemo/broadphase/PairBench.cpp | 184 +++++++++++++++++++++----- 2 files changed, 151 insertions(+), 38 deletions(-) diff --git a/demo/gpudemo/GpuDemo.cpp b/demo/gpudemo/GpuDemo.cpp index 57f542b91..afaef84da 100644 --- a/demo/gpudemo/GpuDemo.cpp +++ b/demo/gpudemo/GpuDemo.cpp @@ -45,14 +45,15 @@ void GpuDemo::initCL(int preferredDeviceIndex, int preferredPlatformIndex) //cl_device_type deviceType = CL_DEVICE_TYPE_CPU; //#endif - + cl_platform_id platformId; // 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); + m_clData->m_clContext = btOpenCLUtils::createContextFromType(deviceType, &ciErrNum, 0,0,preferredDeviceIndex, preferredPlatformIndex,&platformId); + btOpenCLUtils::printPlatformInfo(platformId); } diff --git a/demo/gpudemo/broadphase/PairBench.cpp b/demo/gpudemo/broadphase/PairBench.cpp index b989772ee..84f480eb9 100644 --- a/demo/gpudemo/broadphase/PairBench.cpp +++ b/demo/gpudemo/broadphase/PairBench.cpp @@ -23,13 +23,73 @@ __kernel void moveObjectsKernel(__global float4* posOrnColors, int numObjects) if (iGID>=numObjects) return; __global float4* positions = &posOrnColors[0]; - positions[iGID].y +=0.01f; + if (iGID<0.5*numObjects) + { + positions[iGID].y +=0.01f; + } __global float4* colors = &posOrnColors[numObjects*2]; - float fi = iGID; - float c = fi/numObjects; - colors[iGID] = (float4)(0,0,c,1); - + colors[iGID] = (float4)(0,0,1,1); } + +__kernel void colorPairsKernel(__global float4* posOrnColors, int numObjects, __global const int2* pairs, int numPairs) +{ + int iPairId = get_global_id(0); + if (iPairId>=numPairs) + return; + __global float4* colors = &posOrnColors[numObjects*2]; + + int iObjectA = pairs[iPairId].x; + int iObjectB = pairs[iPairId].y; + colors[iObjectA] = (float4)(1,0,0,1); + colors[iObjectB] = (float4)(1,0,0,1); +} + +__kernel void + sineWaveKernel( __global float4* posOrnColors, __global float* pBodyTimes,const int numNodes) +{ + int nodeID = get_global_id(0); + float timeStepPos = 0.000166666; + float mAmplitude = 26.f; + if( nodeID < numNodes ) + { + pBodyTimes[nodeID] += timeStepPos; + float4 position = posOrnColors[nodeID]; + position.x = native_cos(pBodyTimes[nodeID]*2.17f)*mAmplitude + native_sin(pBodyTimes[nodeID])*mAmplitude*0.5f; + position.y = native_cos(pBodyTimes[nodeID]*1.38f)*mAmplitude + native_sin(pBodyTimes[nodeID]*mAmplitude); + position.z = native_cos(pBodyTimes[nodeID]*2.17f)*mAmplitude + native_sin(pBodyTimes[nodeID]*0.777f)*mAmplitude; + + posOrnColors[nodeID] = position; + __global float4* colors = &posOrnColors[numNodes*2]; + colors[nodeID] = (float4)(0,0,1,1); + } +} + +typedef struct +{ + float fx; + float fy; + float fz; + int uw; +} btAABBCL; + +__kernel void updateAabbSimple( __global float4* posOrnColors, const int numNodes, __global btAABBCL* pAABB) +{ + int nodeId = get_global_id(0); + if( nodeId < numNodes ) + { + float4 position = posOrnColors[nodeId]; + float4 halfExtents = (float4)(1.01f,1.01f,1.01f,0.f); + pAABB[nodeId*2].fx = position.x-halfExtents.x; + pAABB[nodeId*2].fy = position.y-halfExtents.y; + pAABB[nodeId*2].fz = position.z-halfExtents.z; + pAABB[nodeId*2].uw = nodeId; + pAABB[nodeId*2+1].fx = position.x+halfExtents.x; + pAABB[nodeId*2+1].fy = position.y+halfExtents.y; + pAABB[nodeId*2+1].fz = position.z+halfExtents.z; + pAABB[nodeId*2+1].uw = nodeId; + } +} + ); @@ -38,17 +98,20 @@ struct PairBenchInternalData btGpuSapBroadphase* m_broadphaseGPU; cl_kernel m_moveObjectsKernel; - cl_kernel m_colorObjectsKernel; + cl_kernel m_sineWaveKernel; cl_kernel m_colorPairsKernel; + cl_kernel m_updateAabbSimple; btOpenCLArray* m_instancePosOrnColor; - + btOpenCLArray* m_bodyTimes; PairBenchInternalData() :m_broadphaseGPU(0), m_moveObjectsKernel(0), - m_colorObjectsKernel(0), + m_sineWaveKernel(0), m_colorPairsKernel(0), - m_instancePosOrnColor(0) + m_instancePosOrnColor(0), + m_bodyTimes(0), + m_updateAabbSimple(0) { } }; @@ -93,7 +156,10 @@ void PairBench::initPhysics(const ConstructionInfo& ci) cl_program pairBenchProg=0; int errNum=0; m_data->m_moveObjectsKernel = btOpenCLUtils::compileCLKernelFromString(m_clData->m_clContext,m_clData->m_clDevice,s_pairBenchKernelString,"moveObjectsKernel",&errNum,pairBenchProg); - + m_data->m_sineWaveKernel = btOpenCLUtils::compileCLKernelFromString(m_clData->m_clContext,m_clData->m_clDevice,s_pairBenchKernelString,"sineWaveKernel",&errNum,pairBenchProg); + m_data->m_colorPairsKernel = btOpenCLUtils::compileCLKernelFromString(m_clData->m_clContext,m_clData->m_clDevice,s_pairBenchKernelString,"colorPairsKernel",&errNum,pairBenchProg); + m_data->m_updateAabbSimple = btOpenCLUtils::compileCLKernelFromString(m_clData->m_clContext,m_clData->m_clDevice,s_pairBenchKernelString,"updateAabbSimple",&errNum,pairBenchProg); + } if (ci.m_window) @@ -121,10 +187,10 @@ void PairBench::initPhysics(const ConstructionInfo& ci) { for (int k=0;kregisterGraphicsInstance(shapeId,position,orn,color,scaling); btVector3 aabbHalfExtents(1,1,1); @@ -141,7 +207,7 @@ void PairBench::initPhysics(const ConstructionInfo& ci) float camPos[4]={15.5,12.5,15.5,0}; m_instancingRenderer->setCameraTargetPosition(camPos); - m_instancingRenderer->setCameraDistance(40); + m_instancingRenderer->setCameraDistance(60); m_instancingRenderer->writeTransforms(); m_data->m_broadphaseGPU->writeAabbsToGpu(); @@ -151,6 +217,7 @@ void PairBench::exitPhysics() { delete m_data->m_broadphaseGPU; delete m_data->m_instancePosOrnColor; + delete m_data->m_bodyTimes; m_data->m_broadphaseGPU = 0; m_window->setKeyboardCallback(oldCallback); @@ -168,10 +235,12 @@ void PairBench::clientMoveAndDisplay() //color all objects blue bool animate=true; + int numObjects= m_instancingRenderer->getInternalData()->m_totalNumInstances; + btVector4* positions = 0; if (animate) { GLuint vbo = m_instancingRenderer->getInternalData()->m_vbo; - int numObjects= m_instancingRenderer->getInternalData()->m_totalNumInstances; + int arraySizeInBytes = numObjects * (3)*sizeof(btVector4); @@ -181,8 +250,8 @@ void PairBench::clientMoveAndDisplay() char* hostPtr= (char*)glMapBufferRange( GL_ARRAY_BUFFER,m_instancingRenderer->getMaxShapeCapacity(),arraySizeInBytes, GL_MAP_WRITE_BIT|GL_MAP_READ_BIT );//GL_READ_WRITE);//GL_WRITE_ONLY GLint err = glGetError(); assert(err==GL_NO_ERROR); - - btVector4* positions = (btVector4*)hostPtr; + positions = (btVector4*)hostPtr; + if (m_data->m_instancePosOrnColor && m_data->m_instancePosOrnColor->size() != 3*numObjects) { delete m_data->m_instancePosOrnColor; @@ -193,37 +262,55 @@ void PairBench::clientMoveAndDisplay() m_data->m_instancePosOrnColor = new btOpenCLArray(m_clData->m_clContext,m_clData->m_clQueue); m_data->m_instancePosOrnColor->resize(3*numObjects); m_data->m_instancePosOrnColor->copyFromHostPointer(positions,3*numObjects,0); + m_data->m_bodyTimes = new btOpenCLArray(m_clData->m_clContext,m_clData->m_clQueue); + m_data->m_bodyTimes ->resize(numObjects); + btAlignedObjectArray tmp; + tmp.resize(numObjects); + for (int i=0;im_bodyTimes->copyFromHost(tmp); } - bool animateHost = false; - if (animateHost) + if (1) { - for (int i=0;im_clQueue, m_data->m_moveObjectsKernel); - - launcher.setBuffer(m_data->m_instancePosOrnColor->getBufferCL() ); - launcher.setConst( numObjects); - launcher.launch1D( numObjects); - clFinish(m_clData->m_clQueue); - m_data->m_instancePosOrnColor->copyToHostPointer(positions,3*numObjects,0); + btLauncherCL launcher(m_clData->m_clQueue, m_data->m_sineWaveKernel); + launcher.setBuffer(m_data->m_instancePosOrnColor->getBufferCL() ); + launcher.setBuffer(m_data->m_bodyTimes->getBufferCL() ); + launcher.setConst( numObjects); + launcher.launch1D( numObjects); + clFinish(m_clData->m_clQueue); + } + else + { + + btLauncherCL launcher(m_clData->m_clQueue, m_data->m_moveObjectsKernel); + launcher.setBuffer(m_data->m_instancePosOrnColor->getBufferCL() ); + launcher.setConst( numObjects); + launcher.launch1D( numObjects); + clFinish(m_clData->m_clQueue); + } } } + { + btLauncherCL launcher(m_clData->m_clQueue, m_data->m_updateAabbSimple); + launcher.setBuffer(m_data->m_instancePosOrnColor->getBufferCL() ); + launcher.setConst( numObjects); + launcher.setBuffer(m_data->m_broadphaseGPU->getAabbBuffer()); + launcher.launch1D( numObjects); + clFinish(m_clData->m_clQueue); + + } { BT_PROFILE("calculateOverlappingPairs"); m_data->m_broadphaseGPU->calculateOverlappingPairs(); - int numPairs = m_data->m_broadphaseGPU->getNumOverlap(); + //int numPairs = m_data->m_broadphaseGPU->getNumOverlap(); + //printf("numPairs = %d\n", numPairs); } if (animate) @@ -231,6 +318,31 @@ void PairBench::clientMoveAndDisplay() GLint err = glGetError(); assert(err==GL_NO_ERROR); //color overlapping objects in red + + + if (m_data->m_broadphaseGPU->getNumOverlap()) + { + bool colorPairsOnHost = false; + if (colorPairsOnHost ) + { + + } else + { + int numPairs = m_data->m_broadphaseGPU->getNumOverlap(); + cl_mem pairBuf = m_data->m_broadphaseGPU->getOverlappingPairBuffer(); + + btLauncherCL launcher(m_clData->m_clQueue, m_data->m_colorPairsKernel); + launcher.setBuffer(m_data->m_instancePosOrnColor->getBufferCL() ); + launcher.setConst( numObjects); + launcher.setBuffer( pairBuf); + launcher.setConst( numPairs); + launcher.launch1D( numPairs); + clFinish(m_clData->m_clQueue); + } + } + + m_data->m_instancePosOrnColor->copyToHostPointer(positions,3*numObjects,0); + glUnmapBuffer( GL_ARRAY_BUFFER); err = glGetError(); assert(err==GL_NO_ERROR);