Bart/bullet3
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

improved PairBench.

Browse Source
This commit is contained in:
erwin coumans
2013-03-14 14:35:19 -07:00
parent 5784d6ce0b
commit 39dbb51f68
2 changed files with 151 additions and 38 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
demo/gpudemo/GpuDemo.cpp
View File

@@ -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);
}

184
demo/gpudemo/broadphase/PairBench.cpp
View File

@@ -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<btVector4>* m_instancePosOrnColor;
btOpenCLArray<float>* 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;k<ci.arraySizeZ;k++)
{
btVector3 position(k*2,i*2,j*2);
btVector3 position(k*3,i*3,j*3);
btQuaternion orn(1,0,0,0);
btVector4 color(0,0,1,1);
btVector4 color(0,1,0,1);
btVector4 scaling(1,1,1,1);
int id = ci.m_instancingRenderer->registerGraphicsInstance(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<btVector4>(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<float>(m_clData->m_clContext,m_clData->m_clQueue);
m_data->m_bodyTimes ->resize(numObjects);
btAlignedObjectArray<float> tmp;
tmp.resize(numObjects);
for (int i=0;i<numObjects;i++)
{
tmp[i] = float(i)*(1024.f/numObjects);
}
m_data->m_bodyTimes->copyFromHost(tmp);
}
bool animateHost = false;
if (animateHost)
if (1)
{
for (int i=0;i<numObjects;i++)
if (1)
{
positions[i][1]+=0.1f;
}
btVector4* colors= (btVector4*)(hostPtr+8*sizeof(float)*numObjects);
for (int i=0;i<numObjects;i++)
{
colors[i].setValue(0,0,float(i)/float(numObjects),1);
}
} 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);
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);