From 7cdda65822cae8ba6baed424414202e10ae9510d Mon Sep 17 00:00:00 2001 From: erwincoumans Date: Sun, 26 Jan 2014 11:27:20 -0800 Subject: [PATCH] Experimenting with GUI in PairBench, still preliminary. Goal is to make it easier to add variables that can be tuned using GUI/gwen Pre-compile PairBench OpenCL kernel --- Demos3/GpuDemos/broadphase/PairBench.cpp | 373 +++++++++++++----- Demos3/GpuDemos/broadphase/PairBench.h | 7 +- Demos3/GpuDemos/broadphase/pairsKernel.cl | 70 ++++ Demos3/GpuDemos/broadphase/pairsKernel.h | 68 ++++ Demos3/GpuDemos/gwenInternalData.h | 51 +++ Demos3/GpuDemos/gwenUserInterface.cpp | 60 +-- Demos3/GpuDemos/gwenUserInterface.h | 9 + build3/stringify.bat | 1 + .../b3GpuBroadphaseInterface.h | 1 + .../b3GpuGridBroadphase.cpp | 66 ++-- .../BroadphaseCollision/b3GpuGridBroadphase.h | 6 + .../b3GpuSapBroadphase.cpp | 2 + .../BroadphaseCollision/b3GpuSapBroadphase.h | 9 +- .../kernels/gridBroadphase.cl | 190 +-------- .../kernels/gridBroadphaseKernels.h | 188 +-------- 15 files changed, 567 insertions(+), 534 deletions(-) create mode 100644 Demos3/GpuDemos/broadphase/pairsKernel.cl create mode 100644 Demos3/GpuDemos/broadphase/pairsKernel.h create mode 100644 Demos3/GpuDemos/gwenInternalData.h diff --git a/Demos3/GpuDemos/broadphase/PairBench.cpp b/Demos3/GpuDemos/broadphase/PairBench.cpp index abc5edc94..c56a93551 100644 --- a/Demos3/GpuDemos/broadphase/PairBench.cpp +++ b/Demos3/GpuDemos/broadphase/PairBench.cpp @@ -13,94 +13,45 @@ #include "Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h" #include "../../../btgui/Timing/b3Quickprof.h" #include "../gwenUserInterface.h" +#include "../GwenInternalData.h" + #include +#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 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;ideleteBroadphase(); + sPairDemo->createBroadphase(dimensions[0],dimensions[1],dimensions[2]); + break; + } + } + } + + } + +}; + + +template +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;iAddItem(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* handler = new MySliderEventHandler(labelName,label,&dimensions[i]); + pSlider->onValueChanged.Add( handler, &MySliderEventHandler::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* handler = new MySliderEventHandler(labelName,label,&mAmplitude); + pSlider->onValueChanged.Add( handler, &MySliderEventHandler::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;iregisterGraphicsInstance(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;im_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 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; } diff --git a/Demos3/GpuDemos/broadphase/PairBench.h b/Demos3/GpuDemos/broadphase/PairBench.h index d4380a19e..724ac9130 100644 --- a/Demos3/GpuDemos/broadphase/PairBench.h +++ b/Demos3/GpuDemos/broadphase/PairBench.h @@ -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"; diff --git a/Demos3/GpuDemos/broadphase/pairsKernel.cl b/Demos3/GpuDemos/broadphase/pairsKernel.cl new file mode 100644 index 000000000..3c3b0bc3f --- /dev/null +++ b/Demos3/GpuDemos/broadphase/pairsKernel.cl @@ -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; + } +} diff --git a/Demos3/GpuDemos/broadphase/pairsKernel.h b/Demos3/GpuDemos/broadphase/pairsKernel.h new file mode 100644 index 000000000..d707edf3a --- /dev/null +++ b/Demos3/GpuDemos/broadphase/pairsKernel.h @@ -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" +; diff --git a/Demos3/GpuDemos/gwenInternalData.h b/Demos3/GpuDemos/gwenInternalData.h new file mode 100644 index 000000000..5643f2316 --- /dev/null +++ b/Demos3/GpuDemos/gwenInternalData.h @@ -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 m_handlers; + b3ToggleButtonCallback m_toggleButtonCallback; + b3ComboBoxCallback m_comboBoxCallback; + + +}; + +#endif//GWEN_INTERNAL_DATA_H diff --git a/Demos3/GpuDemos/gwenUserInterface.cpp b/Demos3/GpuDemos/gwenUserInterface.cpp index e8c60c8a9..2a9ed96b4 100644 --- a/Demos3/GpuDemos/gwenUserInterface.cpp +++ b/Demos3/GpuDemos/gwenUserInterface.cpp @@ -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 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; + + } diff --git a/Demos3/GpuDemos/gwenUserInterface.h b/Demos3/GpuDemos/gwenUserInterface.h index 443df3e66..4c32376bf 100644 --- a/Demos3/GpuDemos/gwenUserInterface.h +++ b/Demos3/GpuDemos/gwenUserInterface.h @@ -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 diff --git a/build3/stringify.bat b/build3/stringify.bat index dddceef20..7c65b8149 100644 --- a/build3/stringify.bat +++ b/build3/stringify.bat @@ -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 diff --git a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuBroadphaseInterface.h b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuBroadphaseInterface.h index 7e5ba32ee..3598ffb56 100644 --- a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuBroadphaseInterface.h +++ b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuBroadphaseInterface.h @@ -13,6 +13,7 @@ class b3GpuBroadphaseInterface { public: + typedef class b3GpuBroadphaseInterface* (CreateFunc)(cl_context ctx,cl_device_id device, cl_command_queue q); virtual ~b3GpuBroadphaseInterface() { diff --git a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.cpp b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.cpp index 43b17bad6..a54886d9f 100644 --- a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.cpp +++ b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuGridBroadphase.cpp @@ -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 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 pairsGpu2(m_context,m_queue); - b3OpenCLArray pairsGpu(m_context,m_queue); - b3OpenCLArray pairStartCurGpu(m_context,m_queue); - b3AlignedObjectArray 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() m_pairCount; + b3OpenCLArray m_allAabbsGPU; b3AlignedObjectArray 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 \ No newline at end of file diff --git a/src/Bullet3OpenCL/BroadphaseCollision/kernels/gridBroadphase.cl b/src/Bullet3OpenCL/BroadphaseCollision/kernels/gridBroadphase.cl index aa453ca41..d7d71250e 100644 --- a/src/Bullet3OpenCL/BroadphaseCollision/kernels/gridBroadphase.cl +++ b/src/Bullet3OpenCL/BroadphaseCollision/kernels/gridBroadphase.cl @@ -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= 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; -} diff --git a/src/Bullet3OpenCL/BroadphaseCollision/kernels/gridBroadphaseKernels.h b/src/Bullet3OpenCL/BroadphaseCollision/kernels/gridBroadphaseKernels.h index 881d783c0..b1f7db368 100644 --- a/src/Bullet3OpenCL/BroadphaseCollision/kernels/gridBroadphaseKernels.h +++ b/src/Bullet3OpenCL/BroadphaseCollision/kernels/gridBroadphaseKernels.h @@ -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= 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" ;