AuroraMiddleware/bullet3
1
0
Fork 0
mirror of https://github.com/bulletphysics/bullet3 synced 2025-01-18 21:10:05 +00:00
Code Issues Releases Wiki Activity

Experimenting with GUI in PairBench, still preliminary. Goal is to make it easier to add variables that can be tuned using GUI/gwen

Browse Source 
Pre-compile PairBench OpenCL kernel
This commit is contained in:
erwincoumans 2014-01-26 11:27:20 -08:00
parent ec4292ad75
commit 7cdda65822
15 changed files with 567 additions and 534 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

373
Demos3/GpuDemos/broadphase/PairBench.cpp
View File

@ -13,94 +13,45 @@
#include "Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h"
#include "../../../btgui/Timing/b3Quickprof.h"
#include "../gwenUserInterface.h"
#include "../GwenInternalData.h"
#include <string.h>
#include "pairsKernel.h"
static b3KeyboardCallback oldCallback = 0;
char* gPairBenchFileName = 0;
extern bool useShadowMap;
float maxExtents = -1e30f;
int largeCount = 0;
float timeStepPos = 0.000166666;
float mAmplitude = 251.f;
int dimensions[3]={10,10,10};
const char* axisNames[3] = {"# x-axis","# y-axis","# z-axis"};
extern bool gReset;
#define MSTRINGIFY(A) #A
static const char* s_pairBenchKernelString = MSTRINGIFY(
__kernel void moveObjectsKernel(__global float4* posOrnColors, int numObjects)
struct BroadphaseEntry
{
int iGID = get_global_id(0);
if (iGID>=numObjects)
return;
__global float4* positions = &posOrnColors[0];
if (iGID<0.5*numObjects)
{
positions[iGID].y +=0.01f;
}
__global float4* colors = &posOrnColors[numObjects*2];
colors[iGID] = (float4)(0,0,1,1);
}
const char* m_name;
b3GpuBroadphaseInterface::CreateFunc* m_createFunc;
};
__kernel void colorPairsKernel(__global float4* posOrnColors, int numObjects, __global const int4* pairs, int numPairs)
static PairBench* sPairDemo = 0;
#define BP_COMBO_INDEX 123
static int curSelectedBroadphase = 0;
static BroadphaseEntry allBroadphases[]=
{
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 = 51.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;
} b3AABBCL;
__kernel void updateAabbSimple( __global float4* posOrnColors, const int numNodes, __global b3AABBCL* 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;
}
}
);
{"Gpu Grid",b3GpuGridBroadphase::CreateFunc},
{"Gpu 1-Sap",b3GpuSapBroadphase::CreateFunc},
};
struct PairBenchInternalData
@ -126,6 +77,10 @@ struct PairBenchInternalData
m_updateAabbSimple(0)
{
}
int m_oldYposition;
b3AlignedObjectArray<Gwen::Controls::Base*> m_myControls;
};
@ -165,27 +120,215 @@ static inline float parseFloat(const char*& token)
return f;
}
extern bool useShadowMap;
enum PairToggleButtons
{
MY_RESET = 1024,
};
#define PAIRS_CL_PROGRAM_PATH "Demos3/GpuDemos/broadphase/pairsKernel.cl"
struct PairComboBoxHander :public Gwen::Event::Handler
{
int m_buttonId;
int m_active;
PairComboBoxHander (int buttonId)
:m_buttonId(buttonId),
m_active(false)
{
}
void onSelect( Gwen::Controls::Base* pControl )
{
if (m_active)
{
Gwen::Controls::ComboBox* but = (Gwen::Controls::ComboBox*) pControl;
Gwen::String str = Gwen::Utility::UnicodeToString( but->GetSelectedItem()->GetText());
int numItems = sizeof(allBroadphases)/sizeof(BroadphaseEntry);
//find selected item
for (int i=0;i<numItems;i++)
{
if (strcmp(str.c_str(),allBroadphases[i].m_name)==0)
{
curSelectedBroadphase=i;
sPairDemo->deleteBroadphase();
sPairDemo->createBroadphase(dimensions[0],dimensions[1],dimensions[2]);
break;
}
}
}
}
};
template<typename T>
struct MySliderEventHandler : public Gwen::Event::Handler
{
Gwen::Controls::TextBox* m_label;
char m_variableName[1024];
T* m_targetValue;
MySliderEventHandler(const char* varName, Gwen::Controls::TextBox* label, T* target)
:m_label(label),
m_targetValue(target)
{
memcpy(m_variableName,varName,strlen(varName)+1);
}
void SliderMoved( Gwen::Controls::Base* pControl )
{
Gwen::Controls::Slider* pSlider = (Gwen::Controls::Slider*)pControl;
//printf("value = %f\n", pSlider->GetValue());//UnitPrint( Utility::Format( L"Slider Value: %.2f", pSlider->GetValue() ) );
char txt[1024];
T v = T(pSlider->GetValue());
(*m_targetValue) = v;
int val = int(v);//todo: specialize on template type
sprintf(txt,"%s : %d", m_variableName,val);
m_label->SetText(txt);
}
};
void PairBench::initPhysics(const ConstructionInfo& ci)
{
m_instancingRenderer = ci.m_instancingRenderer;
sPairDemo = this;
useShadowMap = false;
m_data->m_gui = ci.m_gui;
//remember the old position in the GUI, to restore at exit
GwenInternalData* data = m_data->m_gui->getInternalData();
m_data->m_oldYposition = data->m_curYposition;
data->m_curYposition+=40;
{
int startItem = 0;
int numBroadphases = sizeof(allBroadphases)/sizeof(BroadphaseEntry);
Gwen::Controls::ComboBox* combobox = new Gwen::Controls::ComboBox(data->m_demoPage->GetPage());
PairComboBoxHander* handler = new PairComboBoxHander(555);
m_data->m_myControls.push_back(combobox);
combobox->onSelection.Add(handler,&PairComboBoxHander::onSelect);
int ypos = data->m_curYposition;
combobox->SetPos(10, ypos );
combobox->SetWidth( 100 );
for (int i=0;i<numBroadphases;i++)
{
Gwen::Controls::MenuItem* item = combobox->AddItem(Gwen::Utility::StringToUnicode(allBroadphases[i].m_name));
if (i==startItem)
combobox->OnItemSelected(item);
}
handler->m_active = true;
data->m_curYposition+=22;
}
if (1)
for (int i=0;i<3;i++)
{
{
Gwen::Controls::TextBox* label = new Gwen::Controls::TextBox(data->m_demoPage->GetPage());
m_data->m_myControls.push_back(label);
label->SetText( "Text Label" );
label->SetPos( 10, 10 + 25 );
label->SetWidth(100);
label->SetPos(10,data->m_curYposition);
data->m_curYposition+=22;
Gwen::Controls::HorizontalSlider* pSlider = new Gwen::Controls::HorizontalSlider( data->m_demoPage->GetPage());
m_data->m_myControls.push_back(pSlider);
pSlider->SetPos( 10, data->m_curYposition );
pSlider->SetSize( 100, 20 );
pSlider->SetRange( 0, 100 );
pSlider->SetValue( dimensions[i] );
char labelName[1024];
sprintf(labelName,"%s",axisNames[0]);
MySliderEventHandler<int>* handler = new MySliderEventHandler<int>(labelName,label,&dimensions[i]);
pSlider->onValueChanged.Add( handler, &MySliderEventHandler<int>::SliderMoved );
handler->SliderMoved(pSlider);
float v = pSlider->GetValue();
data->m_curYposition+=22;
}
}
if (1)
{
{
Gwen::Controls::TextBox* label = new Gwen::Controls::TextBox(data->m_demoPage->GetPage());
m_data->m_myControls.push_back(label);
const char* labelName = "Scale: ";
label->SetText( labelName);
label->SetPos( 10, 10 + 25 );
label->SetWidth(100);
label->SetPos(10,data->m_curYposition);
data->m_curYposition+=22;
Gwen::Controls::HorizontalSlider* pSlider = new Gwen::Controls::HorizontalSlider( data->m_demoPage->GetPage());
m_data->m_myControls.push_back(pSlider);
pSlider->SetPos( 10, data->m_curYposition );
pSlider->SetSize( 100, 20 );
pSlider->SetRange( 0, 300);
pSlider->SetValue( mAmplitude );
MySliderEventHandler<float>* handler = new MySliderEventHandler<float>(labelName,label,&mAmplitude);
pSlider->onValueChanged.Add( handler, &MySliderEventHandler<float>::SliderMoved );
handler->SliderMoved(pSlider);
float v = pSlider->GetValue();
data->m_curYposition+=22;
}
}
//pSlider->onValueChanged.Add( this, &Slider::SliderMoved );
data->m_curYposition+=22;
/*m_data->m_gui->registerToggleButton(MY_RESET,"reset");
sOldCallback = m_data->m_gui->getToggleButtonCallback();
m_data->m_gui->setToggleButtonCallback(PairButtonCallback);
*/
int startItem = 0;
//m_data->m_gui->registerComboBox(BP_COMBO_INDEX,numBroadphases,&mydemonames[0],startItem);
//sOldComboCallback = m_data->m_gui->getComboBoxCallback();
//m_data->m_gui->setComboBoxCallback(PairComboBoxCallback);
initCL(ci.preferredOpenCLDeviceIndex,ci.preferredOpenCLPlatformIndex);
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;
cl_int err;
cl_program pairBenchProg=b3OpenCLUtils::compileCLProgramFromString(m_clData->m_clContext,m_clData->m_clDevice,pairsKernelsCL,&err,"",PAIRS_CL_PROGRAM_PATH);
int errNum=0;
m_data->m_moveObjectsKernel = b3OpenCLUtils::compileCLKernelFromString(m_clData->m_clContext,m_clData->m_clDevice,s_pairBenchKernelString,"moveObjectsKernel",&errNum,pairBenchProg);
m_data->m_sineWaveKernel = b3OpenCLUtils::compileCLKernelFromString(m_clData->m_clContext,m_clData->m_clDevice,s_pairBenchKernelString,"sineWaveKernel",&errNum,pairBenchProg);
m_data->m_colorPairsKernel = b3OpenCLUtils::compileCLKernelFromString(m_clData->m_clContext,m_clData->m_clDevice,s_pairBenchKernelString,"colorPairsKernel",&errNum,pairBenchProg);
m_data->m_updateAabbSimple = b3OpenCLUtils::compileCLKernelFromString(m_clData->m_clContext,m_clData->m_clDevice,s_pairBenchKernelString,"updateAabbSimple",&errNum,pairBenchProg);
m_data->m_moveObjectsKernel = b3OpenCLUtils::compileCLKernelFromString(m_clData->m_clContext,m_clData->m_clDevice,pairsKernelsCL,"moveObjectsKernel",&errNum,pairBenchProg);
m_data->m_sineWaveKernel = b3OpenCLUtils::compileCLKernelFromString(m_clData->m_clContext,m_clData->m_clDevice,pairsKernelsCL,"sineWaveKernel",&errNum,pairBenchProg);
m_data->m_colorPairsKernel = b3OpenCLUtils::compileCLKernelFromString(m_clData->m_clContext,m_clData->m_clDevice,pairsKernelsCL,"colorPairsKernel",&errNum,pairBenchProg);
m_data->m_updateAabbSimple = b3OpenCLUtils::compileCLKernelFromString(m_clData->m_clContext,m_clData->m_clDevice,pairsKernelsCL,"updateAabbSimple",&errNum,pairBenchProg);
}
@ -197,16 +340,26 @@ void PairBench::initPhysics(const ConstructionInfo& ci)
}
m_instancingRenderer = ci.m_instancingRenderer;
#ifndef B3_NO_PROFILE
b3ProfileManager::CleanupMemory();
#endif //B3_NO_PROFILE
createBroadphase(ci.arraySizeX,ci.arraySizeY,ci.arraySizeZ);
}
void PairBench::createBroadphase(int arraySizeX, int arraySizeY, int arraySizeZ)
{
m_data->m_broadphaseGPU = (allBroadphases[curSelectedBroadphase].m_createFunc)(m_clData->m_clContext,m_clData->m_clDevice,m_clData->m_clQueue);
int strideInBytes = 9*sizeof(float);
int numVertices = sizeof(cube_vertices)/strideInBytes;
int numIndices = sizeof(cube_vertices)/sizeof(int);
int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int shapeId = m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
int group=1;
int mask=1;
int index=10;
@ -294,12 +447,12 @@ void PairBench::initPhysics(const ConstructionInfo& ci)
if (l>500)
{
b3Vector4 color=b3MakeVector4(0,1,0,0.1);
int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
int id = m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
m_data->m_broadphaseGPU->createLargeProxy(aabbMin,aabbMax,index,group,mask);
} else
{
b3Vector4 color=b3MakeVector4(1,0,0,1);
int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
int id = m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
m_data->m_broadphaseGPU->createProxy(aabbMin,aabbMax,index,group,mask);
index++;
}
@ -327,18 +480,18 @@ void PairBench::initPhysics(const ConstructionInfo& ci)
}
else
{
for (int i=0;i<ci.arraySizeX;i++)
for (int i=0;i<arraySizeX;i++)
{
for (int j=0;j<ci.arraySizeY;j++)
for (int j=0;j<arraySizeY;j++)
{
for (int k=0;k<ci.arraySizeZ;k++)
for (int k=0;k<arraySizeZ;k++)
{
b3Vector3 position=b3MakeVector3(k*3,i*3,j*3);
b3Quaternion orn(0,0,0,1);
b3Vector4 color=b3MakeVector4(0,1,0,1);
b3Vector4 scaling=b3MakeVector4(1,1,1,1);
int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
int id = m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
b3Vector3 aabbHalfExtents=b3MakeVector3(1,1,1);
b3Vector3 aabbMin = position-aabbHalfExtents;
@ -364,17 +517,37 @@ void PairBench::initPhysics(const ConstructionInfo& ci)
m_instancingRenderer->writeTransforms();
m_data->m_broadphaseGPU->writeAabbsToGpu();
}
void PairBench::deleteBroadphase()
{
delete m_data->m_broadphaseGPU;
m_data->m_broadphaseGPU = 0;
delete m_data->m_instancePosOrnColor;
m_data->m_instancePosOrnColor = 0;
delete m_data->m_bodyTimes;
m_data->m_bodyTimes = 0;
m_data->m_broadphaseGPU = 0;
m_instancingRenderer->removeAllInstances();
}
void PairBench::exitPhysics()
{
delete m_data->m_broadphaseGPU;
delete m_data->m_instancePosOrnColor;
delete m_data->m_bodyTimes;
m_data->m_broadphaseGPU = 0;
m_data->m_gui->getInternalData()->m_curYposition = m_data->m_oldYposition;
for (int i=0;i<m_data->m_myControls.size();i++)
{
delete m_data->m_myControls[i];
}
sPairDemo = 0;
m_window->setKeyboardCallback(oldCallback);
exitCL();
}
@ -434,6 +607,9 @@ void PairBench::clientMoveAndDisplay()
b3LauncherCL launcher(m_clData->m_clQueue, m_data->m_sineWaveKernel,"m_sineWaveKernel");
launcher.setBuffer(m_data->m_instancePosOrnColor->getBufferCL() );
launcher.setBuffer(m_data->m_bodyTimes->getBufferCL() );
launcher.setConst(timeStepPos);
launcher.setConst(mAmplitude);
launcher.setConst( numObjects);
launcher.launch1D( numObjects);
clFinish(m_clData->m_clQueue);
@ -471,7 +647,8 @@ void PairBench::clientMoveAndDisplay()
b3AlignedObjectArray<b3Vector4> posOrnColorsCpu;
m_data->m_instancePosOrnColor->copyToHost(posOrnColorsCpu);
if (m_data->m_instancePosOrnColor)
m_data->m_instancePosOrnColor->copyToHost(posOrnColorsCpu);
@ -502,8 +679,8 @@ void PairBench::clientMoveAndDisplay()
int sz = sizeof(b3Int4)*64*numObjects;
m_data->m_broadphaseGPU->calculateOverlappingPairs(16*numObjects);
//int numPairs = m_data->m_broadphaseGPU->getNumOverlap();
//printf("numPairs = %d\n", numPairs);
int numPairs = m_data->m_broadphaseGPU->getNumOverlap();
printf("numPairs = %d\n", numPairs);
dt = cl.getTimeMicroseconds()-dt;
}

7
Demos3/GpuDemos/broadphase/PairBench.h
View File

@ -11,15 +11,20 @@ class PairBench : public GpuDemo
struct PairBenchInternalData* m_data;
public:
PairBench();
virtual ~PairBench();
virtual void initPhysics(const ConstructionInfo& ci);
virtual void exitPhysics();
void createBroadphase(int xdim, int ydim, int zdim);
void deleteBroadphase();
virtual const char* getName()
{
return "PairBench";

70
Demos3/GpuDemos/broadphase/pairsKernel.cl Normal file
View File

@ -0,0 +1,70 @@
__kernel void moveObjectsKernel(__global float4* posOrnColors, int numObjects)
{
int iGID = get_global_id(0);
if (iGID>=numObjects)
return;
__global float4* positions = &posOrnColors[0];
if (iGID<0.5*numObjects)
{
positions[iGID].y +=0.01f;
}
__global float4* colors = &posOrnColors[numObjects*2];
colors[iGID] = (float4)(0,0,1,1);
}
__kernel void colorPairsKernel(__global float4* posOrnColors, int numObjects, __global const int4* 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,float timeStepPos, float mAmplitude,const int numNodes)
{
int nodeID = get_global_id(0);
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;
} b3AABBCL;
__kernel void updateAabbSimple( __global float4* posOrnColors, const int numNodes, __global b3AABBCL* 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;
}
}

68
Demos3/GpuDemos/broadphase/pairsKernel.h Normal file
View File

@ -0,0 +1,68 @@
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
static const char* pairsKernelsCL= \
"__kernel void moveObjectsKernel(__global float4* posOrnColors, int numObjects)\n"
"{\n"
" int iGID = get_global_id(0);\n"
" if (iGID>=numObjects)\n"
" return;\n"
" __global float4* positions = &posOrnColors[0];\n"
" if (iGID<0.5*numObjects)\n"
" {\n"
" positions[iGID].y +=0.01f;\n"
" }\n"
" __global float4* colors = &posOrnColors[numObjects*2];\n"
" colors[iGID] = (float4)(0,0,1,1);\n"
"}\n"
"__kernel void colorPairsKernel(__global float4* posOrnColors, int numObjects, __global const int4* pairs, int numPairs)\n"
"{\n"
" int iPairId = get_global_id(0);\n"
" if (iPairId>=numPairs)\n"
" return;\n"
" __global float4* colors = &posOrnColors[numObjects*2];\n"
" int iObjectA = pairs[iPairId].x;\n"
" int iObjectB = pairs[iPairId].y;\n"
" colors[iObjectA] = (float4)(1,0,0,1);\n"
" colors[iObjectB] = (float4)(1,0,0,1);\n"
"}\n"
"__kernel void \n"
" sineWaveKernel( __global float4* posOrnColors, __global float* pBodyTimes,float timeStepPos, float mAmplitude,const int numNodes)\n"
"{\n"
" int nodeID = get_global_id(0);\n"
" if( nodeID < numNodes )\n"
" {\n"
" pBodyTimes[nodeID] += timeStepPos;\n"
" float4 position = posOrnColors[nodeID];\n"
" position.x = native_cos(pBodyTimes[nodeID]*2.17f)*mAmplitude + native_sin(pBodyTimes[nodeID])*mAmplitude*0.5f;\n"
" position.y = native_cos(pBodyTimes[nodeID]*1.38f)*mAmplitude + native_sin(pBodyTimes[nodeID]*mAmplitude);\n"
" position.z = native_cos(pBodyTimes[nodeID]*2.17f)*mAmplitude + native_sin(pBodyTimes[nodeID]*0.777f)*mAmplitude;\n"
" \n"
" posOrnColors[nodeID] = position;\n"
" __global float4* colors = &posOrnColors[numNodes*2];\n"
" colors[nodeID] = (float4)(0,0,1,1);\n"
" }\n"
"}\n"
"typedef struct \n"
"{\n"
" float fx;\n"
" float fy;\n"
" float fz;\n"
" int uw;\n"
"} b3AABBCL;\n"
"__kernel void updateAabbSimple( __global float4* posOrnColors, const int numNodes, __global b3AABBCL* pAABB)\n"
"{\n"
" int nodeId = get_global_id(0);\n"
" if( nodeId < numNodes )\n"
" {\n"
" float4 position = posOrnColors[nodeId];\n"
" float4 halfExtents = (float4)(1.01f,1.01f,1.01f,0.f);\n"
" pAABB[nodeId*2].fx = position.x-halfExtents.x;\n"
" pAABB[nodeId*2].fy = position.y-halfExtents.y;\n"
" pAABB[nodeId*2].fz = position.z-halfExtents.z;\n"
" pAABB[nodeId*2].uw = nodeId;\n"
" pAABB[nodeId*2+1].fx = position.x+halfExtents.x;\n"
" pAABB[nodeId*2+1].fy = position.y+halfExtents.y;\n"
" pAABB[nodeId*2+1].fz = position.z+halfExtents.z;\n"
" pAABB[nodeId*2+1].uw = nodeId; \n"
" }\n"
"}\n"
;

51
Demos3/GpuDemos/gwenInternalData.h Normal file
View File

@ -0,0 +1,51 @@
#ifndef GWEN_INTERNAL_DATA_H
#define GWEN_INTERNAL_DATA_H
#include "OpenGLWindow/GwenOpenGL3CoreRenderer.h"
#include "OpenGLWindow/GLPrimitiveRenderer.h"
#include "Gwen/Platform.h"
#include "Gwen/Controls/TreeControl.h"
#include "Gwen/Controls/RadioButtonController.h"
#include "Gwen/Controls/VerticalSlider.h"
#include "Gwen/Controls/HorizontalSlider.h"
#include "Gwen/Controls/GroupBox.h"
#include "Gwen/Controls/CheckBox.h"
#include "Gwen/Controls/StatusBar.h"
#include "Gwen/Controls/Button.h"
#include "Gwen/Controls/ComboBox.h"
#include "Gwen/Controls/MenuStrip.h"
#include "Gwen/Controls/Slider.h"
#include "Gwen/Controls/Property/Text.h"
#include "Gwen/Controls/SplitterBar.h"
#include "Bullet3Common/b3AlignedObjectArray.h"
#include "Gwen/Gwen.h"
#include "Gwen/Align.h"
#include "Gwen/Utility.h"
#include "Gwen/Controls/WindowControl.h"
#include "Gwen/Controls/TabControl.h"
#include "Gwen/Controls/ListBox.h"
#include "Gwen/Skins/Simple.h"
//#include "Gwen/Skins/TexturedBase.h"
struct GwenInternalData
{
struct sth_stash;
class GwenOpenGL3CoreRenderer* pRenderer;
Gwen::Skin::Simple skin;
Gwen::Controls::Canvas* pCanvas;
GLPrimitiveRenderer* m_primRenderer;
Gwen::Controls::TabButton* m_demoPage;
int m_curYposition;
Gwen::Controls::Label* m_rightStatusBar;
Gwen::Controls::Label* m_leftStatusBar;
b3AlignedObjectArray<struct Gwen::Event::Handler*> m_handlers;
b3ToggleButtonCallback m_toggleButtonCallback;
b3ComboBoxCallback m_comboBoxCallback;
};
#endif//GWEN_INTERNAL_DATA_H

60
Demos3/GpuDemos/gwenUserInterface.cpp
View File

@ -1,48 +1,7 @@
#include "gwenUserInterface.h"
#include "OpenGLWindow/GwenOpenGL3CoreRenderer.h"
#include "OpenGLWindow/GLPrimitiveRenderer.h"
#include "Gwen/Platform.h"
#include "Gwen/Controls/TreeControl.h"
#include "Gwen/Controls/RadioButtonController.h"
#include "Gwen/Controls/VerticalSlider.h"
#include "Gwen/Controls/HorizontalSlider.h"
#include "Gwen/Controls/GroupBox.h"
#include "Gwen/Controls/CheckBox.h"
#include "Gwen/Controls/StatusBar.h"
#include "Gwen/Controls/Button.h"
#include "Gwen/Controls/ComboBox.h"
#include "Gwen/Controls/MenuStrip.h"
#include "Gwen/Controls/Property/Text.h"
#include "Gwen/Controls/SplitterBar.h"
#include "Bullet3Common/b3AlignedObjectArray.h"
#include "Gwen/Gwen.h"
#include "Gwen/Align.h"
#include "Gwen/Utility.h"
#include "Gwen/Controls/WindowControl.h"
#include "Gwen/Controls/TabControl.h"
#include "Gwen/Controls/ListBox.h"
#include "Gwen/Skins/Simple.h"
//#include "Gwen/Skins/TexturedBase.h"
#include "gwenInternalData.h"
struct GwenInternalData
{
struct sth_stash;
class GwenOpenGL3CoreRenderer* pRenderer;
Gwen::Skin::Simple skin;
Gwen::Controls::Canvas* pCanvas;
GLPrimitiveRenderer* m_primRenderer;
Gwen::Controls::TabButton* m_demoPage;
Gwen::Controls::Label* m_rightStatusBar;
Gwen::Controls::Label* m_leftStatusBar;
b3AlignedObjectArray<struct Gwen::Event::Handler*> m_handlers;
b3ToggleButtonCallback m_toggleButtonCallback;
b3ComboBoxCallback m_comboBoxCallback;
};
GwenUserInterface::GwenUserInterface()
{
m_data = new GwenInternalData();
@ -170,6 +129,7 @@ void GwenUserInterface::setStatusBarMessage(const char* message, bool isLeft)
void GwenUserInterface::init(int width, int height,struct sth_stash* stash,float retinaScale)
{
m_data->m_curYposition = 20;
m_data->m_primRenderer = new GLPrimitiveRenderer(width,height);
m_data->pRenderer = new GwenOpenGL3CoreRenderer(m_data->m_primRenderer,stash,width,height,retinaScale);
@ -259,6 +219,10 @@ void GwenUserInterface::init(int width, int height,struct sth_stash* stash,float
*/
}
b3ToggleButtonCallback GwenUserInterface::getToggleButtonCallback()
{
return m_data->m_toggleButtonCallback;
}
void GwenUserInterface::setToggleButtonCallback(b3ToggleButtonCallback callback)
{
@ -272,13 +236,14 @@ void GwenUserInterface::registerToggleButton(int buttonId, const char* name)
Gwen::Controls::Button* but = new Gwen::Controls::Button(m_data->m_demoPage->GetPage());
///some heuristic to find the button location
int ypos = m_data->m_handlers.size()*20;
int ypos = m_data->m_curYposition;
but->SetPos(10, ypos );
but->SetWidth( 100 );
//but->SetBounds( 200, 30, 300, 200 );
MyButtonHander* handler = new MyButtonHander(m_data, buttonId);
m_data->m_handlers.push_back(handler);
m_data->m_curYposition+=22;
but->onToggle.Add(handler, &MyButtonHander::onButtonA);
but->SetIsToggle(true);
but->SetToggleState(false);
@ -291,6 +256,10 @@ void GwenUserInterface::setComboBoxCallback(b3ComboBoxCallback callback)
m_data->m_comboBoxCallback = callback;
}
b3ComboBoxCallback GwenUserInterface::getComboBoxCallback()
{
return m_data->m_comboBoxCallback;
}
void GwenUserInterface::registerComboBox(int comboboxId, int numItems, const char** items, int startItem)
{
Gwen::Controls::ComboBox* combobox = new Gwen::Controls::ComboBox(m_data->m_demoPage->GetPage());
@ -298,7 +267,7 @@ void GwenUserInterface::registerComboBox(int comboboxId, int numItems, const cha
m_data->m_handlers.push_back(handler);
combobox->onSelection.Add(handler,&MyComboBoxHander::onSelect);
int ypos = m_data->m_handlers.size()*20;
int ypos = m_data->m_curYposition;
combobox->SetPos(10, ypos );
combobox->SetWidth( 100 );
//box->SetPos(120,130);
@ -309,6 +278,9 @@ void GwenUserInterface::registerComboBox(int comboboxId, int numItems, const cha
combobox->OnItemSelected(item);
}
m_data->m_curYposition+=22;
}

9
Demos3/GpuDemos/gwenUserInterface.h
View File

@ -26,12 +26,21 @@ class GwenUserInterface
bool mouseButtonCallback(int button, int state, float x, float y);
void setToggleButtonCallback(b3ToggleButtonCallback callback);
b3ToggleButtonCallback getToggleButtonCallback();
void registerToggleButton(int buttonId, const char* name);
void setComboBoxCallback(b3ComboBoxCallback callback);
b3ComboBoxCallback getComboBoxCallback();
void registerComboBox(int buttonId, int numItems, const char** items, int startItem = 0);
void setStatusBarMessage(const char* message, bool isLeft=true);
GwenInternalData* getInternalData()
{
return m_data;
}
};
#endif //_GWEN_USER_INTERFACE_H

1
build3/stringify.bat
View File

@ -58,6 +58,7 @@ premake4 --file=stringifyKernel.lua --kernelfile="../btgui/OpenGLWindow/Shaders/
premake4 --file=stringifyKernel.lua --kernelfile="../Demos3/GpuDemos/broadphase/pairsKernel.cl" --headerfile="../Demos3/GpuDemos/broadphase/pairsKernel.h" --stringname="pairsKernelsCL" stringify

1
src/Bullet3OpenCL/BroadphaseCollision/b3GpuBroadphaseInterface.h
View File

@ -13,6 +13,7 @@ class b3GpuBroadphaseInterface
{
public:
typedef class b3GpuBroadphaseInterface* (CreateFunc)(cl_context ctx,cl_device_id device, cl_command_queue q);
virtual ~b3GpuBroadphaseInterface()
{

66
src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.cpp
View File

@ -20,12 +20,12 @@ cl_kernel kFindCellStart;
cl_kernel kFindOverlappingPairs;
cl_kernel m_copyAabbsKernel;
cl_kernel m_sap2Kernel;
cl_kernel kFindPairsLarge;
cl_kernel kComputePairCacheChanges;
cl_kernel kSqueezeOverlappingPairBuff;
int maxPairsPerBody = 64;
//int maxPairsPerBody = 64;
int maxBodiesPerCell = 256;//??
b3GpuGridBroadphase::b3GpuGridBroadphase(cl_context ctx,cl_device_id device, cl_command_queue q )
@ -86,14 +86,9 @@ m_cellStartGpu(ctx,q)
kFindOverlappingPairs = b3OpenCLUtils::compileCLKernelFromString(m_context, m_device,gridBroadphaseCL, "kFindOverlappingPairs",&errNum,gridProg);
b3Assert(errNum==CL_SUCCESS);
kFindPairsLarge = b3OpenCLUtils::compileCLKernelFromString(m_context, m_device,gridBroadphaseCL, "kFindPairsLarge",&errNum,gridProg);
b3Assert(errNum==CL_SUCCESS);
kComputePairCacheChanges = b3OpenCLUtils::compileCLKernelFromString(m_context, m_device,gridBroadphaseCL, "kComputePairCacheChanges",&errNum,gridProg);
b3Assert(errNum==CL_SUCCESS);
kSqueezeOverlappingPairBuff = b3OpenCLUtils::compileCLKernelFromString(m_context, m_device,gridBroadphaseCL, "kSqueezeOverlappingPairBuff",&errNum,gridProg);
b3Assert(errNum==CL_SUCCESS);
}
m_sorter = new b3RadixSort32CL(m_context,m_device,m_queue);
@ -107,9 +102,9 @@ b3GpuGridBroadphase::~b3GpuGridBroadphase()
clReleaseKernel( kFindOverlappingPairs);
clReleaseKernel( m_sap2Kernel);
clReleaseKernel( m_copyAabbsKernel);
clReleaseKernel( kFindPairsLarge);
clReleaseKernel( kComputePairCacheChanges);
clReleaseKernel( kSqueezeOverlappingPairBuff);
delete m_sorter;
}
@ -201,7 +196,7 @@ void b3GpuGridBroadphase::calculateOverlappingPairs(int maxPairs)
b3OpenCLArray<int> pairCount(m_context,m_queue);
pairCount.push_back(0);
m_gpuPairs.resize(numSmallAabbs*maxPairsPerBody);
m_gpuPairs.resize(maxPairs);//numSmallAabbs*maxPairsPerBody);
{
int numLargeAabbs = m_largeAabbsGPU.size();
@ -281,46 +276,30 @@ void b3GpuGridBroadphase::calculateOverlappingPairs(int maxPairs)
{
B3_PROFILE("kFindOverlappingPairs");
b3OpenCLArray<b3Int2> pairsGpu2(m_context,m_queue);
b3OpenCLArray<unsigned int> pairsGpu(m_context,m_queue);
b3OpenCLArray<unsigned int> pairStartCurGpu(m_context,m_queue);
b3AlignedObjectArray<unsigned int> pairStartCpu;
pairsGpu2.resize(numSmallAabbs*maxPairsPerBody);
pairsGpu.resize(numSmallAabbs*maxPairsPerBody);
pairStartCurGpu.resize(numSmallAabbs*2+2);
pairStartCpu.resize(numSmallAabbs*2+2);
pairStartCpu[0] = 0;
pairStartCpu[1] = 0;
for(int i = 1; i <= numSmallAabbs; i++)
{
pairStartCpu[i * 2] = pairStartCpu[(i-1) * 2] + maxPairsPerBody;
pairStartCpu[i * 2 + 1] = 0;
}
pairStartCurGpu.copyFromHost(pairStartCpu);
b3LauncherCL launch(m_queue,kFindOverlappingPairs,"kFindOverlappingPairs");
launch.setConst(numSmallAabbs);
launch.setBuffer(m_smallAabbsGPU.getBufferCL());
launch.setBuffer(m_hashGpu.getBufferCL());
launch.setBuffer(m_cellStartGpu.getBufferCL());
launch.setBuffer(pairsGpu.getBufferCL());
launch.setBuffer(pairStartCurGpu.getBufferCL());
launch.setBuffer(m_paramsGPU.getBufferCL());
//launch.setBuffer(0);
launch.setBuffer(pairCount.getBufferCL());
launch.setBuffer(m_gpuPairs.getBufferCL());
launch.setConst(maxPairs);
launch.launch1D(numSmallAabbs);
int numPairs = pairCount.at(0);
if (numPairs >maxPairs)
{
b3Error("Error running out of pairs: numPairs = %d, maxPairs = %d.\n", numPairs, maxPairs);
numPairs =maxPairs;
}
int actualCount = pairCount.at(0);
m_gpuPairs.resize(actualCount);
m_gpuPairs.resize(numPairs);
if (0)
{
@ -372,7 +351,10 @@ void b3GpuGridBroadphase::calculateOverlappingPairsHost(int maxPairs)
pair.y = a;//store the original index in the unsorted aabb array
}
m_hostPairs.push_back(pair);
if (m_hostPairs.size()<maxPairs)
{
m_hostPairs.push_back(pair);
}
}
}
}

6
src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.h
View File

@ -54,6 +54,12 @@ public:
b3GpuGridBroadphase(cl_context ctx,cl_device_id device, cl_command_queue q );
virtual ~b3GpuGridBroadphase();
static b3GpuBroadphaseInterface* CreateFunc(cl_context ctx,cl_device_id device, cl_command_queue q)
{
return new b3GpuGridBroadphase(ctx,device,q);
}
virtual void createProxy(const b3Vector3& aabbMin, const b3Vector3& aabbMax, int userPtr ,short int collisionFilterGroup,short int collisionFilterMask);

2
src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.cpp
View File

@ -52,6 +52,8 @@ m_removedCountGPU(ctx,q)
cl_int errNum=0;
b3Assert(m_context);
b3Assert(m_device);
cl_program sapProg = b3OpenCLUtils::compileCLProgramFromString(m_context,m_device,sapSrc,&errNum,"",B3_BROADPHASE_SAP_PATH);
b3Assert(errNum==CL_SUCCESS);
cl_program sapFastProg = b3OpenCLUtils::compileCLProgramFromString(m_context,m_device,sapFastSrc,&errNum,"",B3_BROADPHASE_SAPFAST_PATH);

9
src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.h
View File

@ -52,10 +52,11 @@ class b3GpuSapBroadphase : public b3GpuBroadphaseInterface
int m_currentBuffer;
public:
public:
b3OpenCLArray<int> m_pairCount;
b3OpenCLArray<b3SapAabb> m_allAabbsGPU;
b3AlignedObjectArray<b3SapAabb> m_allAabbsCPU;
@ -89,6 +90,11 @@ class b3GpuSapBroadphase : public b3GpuBroadphaseInterface
b3GpuSapBroadphase(cl_context ctx,cl_device_id device, cl_command_queue q );
virtual ~b3GpuSapBroadphase();
static b3GpuBroadphaseInterface* CreateFunc(cl_context ctx,cl_device_id device, cl_command_queue q)
{
return new b3GpuSapBroadphase(ctx,device,q);
}
virtual void calculateOverlappingPairs(int maxPairs);
virtual void calculateOverlappingPairsHost(int maxPairs);
@ -106,6 +112,7 @@ class b3GpuSapBroadphase : public b3GpuBroadphaseInterface
virtual cl_mem getAabbBufferWS();
virtual int getNumOverlap();
virtual cl_mem getOverlappingPairBuffer();
};
#endif //B3_GPU_SAP_BROADPHASE_H

190
src/Bullet3OpenCL/BroadphaseCollision/kernels/gridBroadphase.cl
View File

@ -103,11 +103,10 @@ void findPairsInCell( int numObjects,
__global int2* pHash,
__global int* pCellStart,
__global float4* pAABB,
__global int* pPairBuff,
__global int2* pPairBuffStartCurr,
__global float4* pParams,
volatile __global int* pairCount,
__global int4* pPairBuff2
__global int4* pPairBuff2,
int maxPairs
)
{
int4 pGridDim = *((__global int4*)(pParams + 1));
@ -125,11 +124,7 @@ void findPairsInCell( int numObjects,
float4 min0 = pAABB[unsorted_indx*2 + 0];
float4 max0 = pAABB[unsorted_indx*2 + 1];
int handleIndex = as_int(min0.w);
int2 start_curr = pPairBuffStartCurr[handleIndex];
int start = start_curr.x;
int curr = start_curr.y;
int2 start_curr_next = pPairBuffStartCurr[handleIndex+1];
int curr_max = start_curr_next.x - start - 1;
int bucketEnd = bucketStart + maxBodiesPerCell;
bucketEnd = (bucketEnd > numObjects) ? numObjects : bucketEnd;
for(int index2 = bucketStart; index2 < bucketEnd; index2++)
@ -153,59 +148,31 @@ void findPairsInCell( int numObjects,
if (handleIndex<handleIndex2)
{
int curPair = atomic_add(pairCount,1);
int4 newpair;
newpair.x = handleIndex;
newpair.y = handleIndex2;
newpair.z = -1;
newpair.w = -1;
pPairBuff2[curPair] = newpair;
if (curPair<maxPairs)
{
int4 newpair;
newpair.x = handleIndex;
newpair.y = handleIndex2;
newpair.z = -1;
newpair.w = -1;
pPairBuff2[curPair] = newpair;
}
}
} else
{
int handleIndex2 = as_int(min1.w);
int k;
for(k = 0; k < curr; k++)
{
int old_pair = pPairBuff[start+k] & (~0x60000000);
if(old_pair == handleIndex2)
{
pPairBuff[start+k] |= 0x40000000;
break;
}
}
if(k == curr)
{
if(curr >= curr_max)
{ // not a good solution, but let's avoid crash
break;
}
pPairBuff[start+curr] = handleIndex2 | 0x20000000;
curr++;
}
}
}
}
}
if (!pairCount)
{
int2 newStartCurr;
newStartCurr.x = start;
newStartCurr.y = curr;
pPairBuffStartCurr[handleIndex] = newStartCurr;
}
}
__kernel void kFindOverlappingPairs( int numObjects,
__global float4* pAABB,
__global int2* pHash,
__global int* pCellStart,
__global int* pPairBuff,
__global int2* pPairBuffStartCurr,
__global float4* pParams ,
volatile __global int* pairCount,
__global int4* pPairBuff2
__global int4* pPairBuff2,
int maxPairs
)
{
@ -235,140 +202,13 @@ __kernel void kFindOverlappingPairs( int numObjects,
for(int x=-1; x<=1; x++)
{
gridPosB.x = gridPosA.x + x;
findPairsInCell(numObjects, gridPosB, index, pHash, pCellStart, pAABB, pPairBuff, pPairBuffStartCurr, pParams, pairCount,pPairBuff2);
findPairsInCell(numObjects, gridPosB, index, pHash, pCellStart, pAABB, pParams, pairCount,pPairBuff2, maxPairs);
}
}
}
}
__kernel void kFindPairsLarge( int numObjects,
__global float4* pAABB,
__global int2* pHash,
__global int* pCellStart,
__global int* pPairBuff,
__global int2* pPairBuffStartCurr,
uint numLarge )
{
int index = get_global_id(0);
if(index >= numObjects)
{
return;
}
int2 sortedData = pHash[index];
int unsorted_indx = sortedData.y;
float4 min0 = pAABB[unsorted_indx*2 + 0];
float4 max0 = pAABB[unsorted_indx*2 + 1];
int handleIndex = as_int(min0.w);
int2 start_curr = pPairBuffStartCurr[handleIndex];
int start = start_curr.x;
int curr = start_curr.y;
int2 start_curr_next = pPairBuffStartCurr[handleIndex+1];
int curr_max = start_curr_next.x - start - 1;
for(uint i = 0; i < numLarge; i++)
{
int indx2 = numObjects + i;
float4 min1 = pAABB[indx2*2 + 0];
float4 max1 = pAABB[indx2*2 + 1];
if(testAABBOverlap(min0, max0, min1, max1))
{
int k;
int handleIndex2 = as_int(min1.w);
for(k = 0; k < curr; k++)
{
int old_pair = pPairBuff[start+k] & (~0x60000000);
if(old_pair == handleIndex2)
{
pPairBuff[start+k] |= 0x40000000;
break;
}
}
if(k == curr)
{
pPairBuff[start+curr] = handleIndex2 | 0x20000000;
if(curr >= curr_max)
{ // not a good solution, but let's avoid crash
break;
}
curr++;
}
}
}
int2 newStartCurr;
newStartCurr.x = start;
newStartCurr.y = curr;
pPairBuffStartCurr[handleIndex] = newStartCurr;
return;
}
__kernel void kComputePairCacheChanges( int numObjects,
__global int* pPairBuff,
__global int2* pPairBuffStartCurr,
__global int* pPairScan,
__global float4* pAABB )
{
int index = get_global_id(0);
if(index >= numObjects)
{
return;
}
float4 bbMin = pAABB[index * 2];
int handleIndex = as_int(bbMin.w);
int2 start_curr = pPairBuffStartCurr[handleIndex];
int start = start_curr.x;
int curr = start_curr.y;
__global int *pInp = pPairBuff + start;
int num_changes = 0;
for(int k = 0; k < curr; k++, pInp++)
{
if(!((*pInp) & 0x40000000))
{
num_changes++;
}
}
pPairScan[index+1] = num_changes;
}
__kernel void kSqueezeOverlappingPairBuff( int numObjects,
__global int* pPairBuff,
__global int2* pPairBuffStartCurr,
__global int* pPairScan,
__global int* pPairOut,
__global float4* pAABB )
{
int index = get_global_id(0);
if(index >= numObjects)
{
return;
}
float4 bbMin = pAABB[index * 2];
int handleIndex = as_int(bbMin.w);
int2 start_curr = pPairBuffStartCurr[handleIndex];
int start = start_curr.x;
int curr = start_curr.y;
__global int* pInp = pPairBuff + start;
__global int* pOut = pPairOut + pPairScan[index+1];
__global int* pOut2 = pInp;
int num = 0;
for(int k = 0; k < curr; k++, pInp++)
{
if(!((*pInp) & 0x40000000))
{
*pOut = *pInp;
pOut++;
}
if((*pInp) & 0x60000000)
{
*pOut2 = (*pInp) & (~0x60000000);
pOut2++;
num++;
}
}
int2 newStartCurr;
newStartCurr.x = start;
newStartCurr.y = num;
pPairBuffStartCurr[handleIndex] = newStartCurr;
}

188
src/Bullet3OpenCL/BroadphaseCollision/kernels/gridBroadphaseKernels.h
View File

@ -92,11 +92,10 @@ static const char* gridBroadphaseCL= \
" __global int2* pHash,\n"
" __global int* pCellStart,\n"
" __global float4* pAABB, \n"
" __global int* pPairBuff,\n"
" __global int2* pPairBuffStartCurr,\n"
" __global float4* pParams,\n"
" volatile __global int* pairCount,\n"
" __global int4* pPairBuff2\n"
" __global int4* pPairBuff2,\n"
" int maxPairs\n"
" )\n"
"{\n"
" int4 pGridDim = *((__global int4*)(pParams + 1));\n"
@ -114,11 +113,7 @@ static const char* gridBroadphaseCL= \
" float4 min0 = pAABB[unsorted_indx*2 + 0]; \n"
" float4 max0 = pAABB[unsorted_indx*2 + 1];\n"
" int handleIndex = as_int(min0.w);\n"
" int2 start_curr = pPairBuffStartCurr[handleIndex];\n"
" int start = start_curr.x;\n"
" int curr = start_curr.y;\n"
" int2 start_curr_next = pPairBuffStartCurr[handleIndex+1];\n"
" int curr_max = start_curr_next.x - start - 1;\n"
" \n"
" int bucketEnd = bucketStart + maxBodiesPerCell;\n"
" bucketEnd = (bucketEnd > numObjects) ? numObjects : bucketEnd;\n"
" for(int index2 = bucketStart; index2 < bucketEnd; index2++) \n"
@ -142,58 +137,30 @@ static const char* gridBroadphaseCL= \
" if (handleIndex<handleIndex2)\n"
" {\n"
" int curPair = atomic_add(pairCount,1);\n"
" int4 newpair;\n"
" newpair.x = handleIndex;\n"
" newpair.y = handleIndex2;\n"
" newpair.z = -1;\n"
" newpair.w = -1;\n"
" pPairBuff2[curPair] = newpair;\n"
" if (curPair<maxPairs)\n"
" {\n"
" int4 newpair;\n"
" newpair.x = handleIndex;\n"
" newpair.y = handleIndex2;\n"
" newpair.z = -1;\n"
" newpair.w = -1;\n"
" pPairBuff2[curPair] = newpair;\n"
" }\n"
" }\n"
" \n"
" } else\n"
" {\n"
" int handleIndex2 = as_int(min1.w);\n"
" int k;\n"
" for(k = 0; k < curr; k++)\n"
" {\n"
" int old_pair = pPairBuff[start+k] & (~0x60000000);\n"
" if(old_pair == handleIndex2)\n"
" {\n"
" pPairBuff[start+k] |= 0x40000000;\n"
" break;\n"
" }\n"
" }\n"
" if(k == curr)\n"
" {\n"
" if(curr >= curr_max) \n"
" { // not a good solution, but let's avoid crash\n"
" break;\n"
" }\n"
" pPairBuff[start+curr] = handleIndex2 | 0x20000000;\n"
" curr++;\n"
" }\n"
" }\n"
" }\n"
" }\n"
" }\n"
" if (!pairCount)\n"
" {\n"
" int2 newStartCurr;\n"
" newStartCurr.x = start;\n"
" newStartCurr.y = curr;\n"
" pPairBuffStartCurr[handleIndex] = newStartCurr;\n"
" }\n"
" \n"
"}\n"
"__kernel void kFindOverlappingPairs( int numObjects,\n"
" __global float4* pAABB, \n"
" __global int2* pHash, \n"
" __global int* pCellStart, \n"
" __global int* pPairBuff, \n"
" __global int2* pPairBuffStartCurr, \n"
" __global float4* pParams ,\n"
" volatile __global int* pairCount,\n"
" __global int4* pPairBuff2\n"
" __global int4* pPairBuff2,\n"
" int maxPairs\n"
" )\n"
"{\n"
" int index = get_global_id(0);\n"
@ -222,134 +189,9 @@ static const char* gridBroadphaseCL= \
" for(int x=-1; x<=1; x++) \n"
" {\n"
" gridPosB.x = gridPosA.x + x;\n"
" findPairsInCell(numObjects, gridPosB, index, pHash, pCellStart, pAABB, pPairBuff, pPairBuffStartCurr, pParams, pairCount,pPairBuff2);\n"
" findPairsInCell(numObjects, gridPosB, index, pHash, pCellStart, pAABB, pParams, pairCount,pPairBuff2, maxPairs);\n"
" }\n"
" }\n"
" }\n"
"}\n"
"__kernel void kFindPairsLarge( int numObjects, \n"
" __global float4* pAABB, \n"
" __global int2* pHash, \n"
" __global int* pCellStart, \n"
" __global int* pPairBuff, \n"
" __global int2* pPairBuffStartCurr, \n"
" uint numLarge )\n"
"{\n"
" int index = get_global_id(0);\n"
" if(index >= numObjects)\n"
" {\n"
" return;\n"
" }\n"
" int2 sortedData = pHash[index];\n"
" int unsorted_indx = sortedData.y;\n"
" float4 min0 = pAABB[unsorted_indx*2 + 0];\n"
" float4 max0 = pAABB[unsorted_indx*2 + 1];\n"
" int handleIndex = as_int(min0.w);\n"
" int2 start_curr = pPairBuffStartCurr[handleIndex];\n"
" int start = start_curr.x;\n"
" int curr = start_curr.y;\n"
" int2 start_curr_next = pPairBuffStartCurr[handleIndex+1];\n"
" int curr_max = start_curr_next.x - start - 1;\n"
" for(uint i = 0; i < numLarge; i++)\n"
" {\n"
" int indx2 = numObjects + i;\n"
" float4 min1 = pAABB[indx2*2 + 0];\n"
" float4 max1 = pAABB[indx2*2 + 1];\n"
" if(testAABBOverlap(min0, max0, min1, max1))\n"
" {\n"
" int k;\n"
" int handleIndex2 = as_int(min1.w);\n"
" for(k = 0; k < curr; k++)\n"
" {\n"
" int old_pair = pPairBuff[start+k] & (~0x60000000);\n"
" if(old_pair == handleIndex2)\n"
" {\n"
" pPairBuff[start+k] |= 0x40000000;\n"
" break;\n"
" }\n"
" }\n"
" if(k == curr)\n"
" {\n"
" pPairBuff[start+curr] = handleIndex2 | 0x20000000;\n"
" if(curr >= curr_max) \n"
" { // not a good solution, but let's avoid crash\n"
" break;\n"
" }\n"
" curr++;\n"
" }\n"
" }\n"
" }\n"
" int2 newStartCurr;\n"
" newStartCurr.x = start;\n"
" newStartCurr.y = curr;\n"
" pPairBuffStartCurr[handleIndex] = newStartCurr;\n"
" return;\n"
"}\n"
"__kernel void kComputePairCacheChanges( int numObjects,\n"
" __global int* pPairBuff, \n"
" __global int2* pPairBuffStartCurr, \n"
" __global int* pPairScan, \n"
" __global float4* pAABB )\n"
"{\n"
" int index = get_global_id(0);\n"
" if(index >= numObjects)\n"
" {\n"
" return;\n"
" }\n"
" float4 bbMin = pAABB[index * 2];\n"
" int handleIndex = as_int(bbMin.w);\n"
" int2 start_curr = pPairBuffStartCurr[handleIndex];\n"
" int start = start_curr.x;\n"
" int curr = start_curr.y;\n"
" __global int *pInp = pPairBuff + start;\n"
" int num_changes = 0;\n"
" for(int k = 0; k < curr; k++, pInp++)\n"
" {\n"
" if(!((*pInp) & 0x40000000))\n"
" {\n"
" num_changes++;\n"
" }\n"
" }\n"
" pPairScan[index+1] = num_changes;\n"
"} \n"
"__kernel void kSqueezeOverlappingPairBuff( int numObjects,\n"
" __global int* pPairBuff, \n"
" __global int2* pPairBuffStartCurr, \n"
" __global int* pPairScan,\n"
" __global int* pPairOut, \n"
" __global float4* pAABB )\n"
"{\n"
" int index = get_global_id(0);\n"
" if(index >= numObjects)\n"
" {\n"
" return;\n"
" }\n"
" float4 bbMin = pAABB[index * 2];\n"
" int handleIndex = as_int(bbMin.w);\n"
" int2 start_curr = pPairBuffStartCurr[handleIndex];\n"
" int start = start_curr.x;\n"
" int curr = start_curr.y;\n"
" __global int* pInp = pPairBuff + start;\n"
" __global int* pOut = pPairOut + pPairScan[index+1];\n"
" __global int* pOut2 = pInp;\n"
" int num = 0; \n"
" for(int k = 0; k < curr; k++, pInp++)\n"
" {\n"
" if(!((*pInp) & 0x40000000))\n"
" {\n"
" *pOut = *pInp;\n"
" pOut++;\n"
" }\n"
" if((*pInp) & 0x60000000)\n"
" {\n"
" *pOut2 = (*pInp) & (~0x60000000);\n"
" pOut2++;\n"
" num++;\n"
" }\n"
" }\n"
" int2 newStartCurr;\n"
" newStartCurr.x = start;\n"
" newStartCurr.y = num;\n"
" pPairBuffStartCurr[handleIndex] = newStartCurr;\n"
"}\n"
;