diff --git a/Demos3/GpuDemos/constraints/ConstraintsDemo.cpp b/Demos3/GpuDemos/constraints/ConstraintsDemo.cpp index 0cac2dcd5..192c5bbef 100644 --- a/Demos3/GpuDemos/constraints/ConstraintsDemo.cpp +++ b/Demos3/GpuDemos/constraints/ConstraintsDemo.cpp @@ -53,7 +53,7 @@ void GpuConstraintsDemo::setupScene(const ConstructionInfo& ci) m_data->m_rigidBodyPipeline->writeAllInstancesToGpu(); // m_data->m_rigidBodyPipeline->setGravity(b3Vector3(4,-10,0)); - float camPos[4]={ci.arraySizeX,ci.gapY,ci.arraySizeZ,0}; + float camPos[4]={ci.arraySizeX,0.5*ci.arraySizeY*ci.gapY,ci.arraySizeZ,0}; //float camPos[4]={1,12.5,1.5,0}; m_instancingRenderer->setCameraTargetPosition(camPos); @@ -167,8 +167,8 @@ int GpuConstraintsDemo::createDynamicsObjects2(const ConstructionInfo& ci, const } //b3Vector3 position((j&1)+i*2.2,1+j*2.,(j&1)+k*2.2); //b3Vector3 position((-ci.arraySizeX/2*ci.gapX)+i*ci.gapX,1+j*2.,(-ci.arraySizeZ/2*ci.gapZ)+k*ci.gapZ); - b3Vector3 position(-ci.arraySizeY/2*2+1+j*2., - 10+i*ci.gapX, + b3Vector3 position(-ci.arraySizeX/2*2+1+j*2., + 10+i*ci.gapY, (-ci.arraySizeZ/2*ci.gapZ)+k*ci.gapZ); b3Quaternion orn(0,0,0,1); @@ -191,7 +191,7 @@ int GpuConstraintsDemo::createDynamicsObjects2(const ConstructionInfo& ci, const { ///enable next line to force CPU constraint solving //c = new b3Point2PointConstraint(pid,prevBody,b3Vector3(-1.1,0,0),b3Vector3(1.1,0,0)); - float breakingThreshold=14; + float breakingThreshold=44; // c->setBreakingImpulseThreshold(breakingThreshold); b3Vector3 pivotInA(-1.1,0,0); b3Vector3 pivotInB (1.1,0,0); @@ -212,7 +212,7 @@ int GpuConstraintsDemo::createDynamicsObjects2(const ConstructionInfo& ci, const b3Quaternion relTargetAB = frameInA.getRotation()*frameInB.getRotation().inverse(); //c = new b3FixedConstraint(pid,prevBody,frameInA,frameInB); - float breakingThreshold = 15;//37.f; + float breakingThreshold = 45;//37.f; //c->setBreakingImpulseThreshold(37.1); int cid = m_data->m_rigidBodyPipeline->createFixedConstraint(pid,prevBody,pivotInA,pivotInB,relTargetAB,breakingThreshold); diff --git a/Demos3/GpuDemos/main_opengl3core.cpp b/Demos3/GpuDemos/main_opengl3core.cpp index c8b607e31..88aae9ca0 100644 --- a/Demos3/GpuDemos/main_opengl3core.cpp +++ b/Demos3/GpuDemos/main_opengl3core.cpp @@ -44,6 +44,7 @@ #include "constraints/ConstraintsDemo.h" bool exportFrame=false; +bool exportMovie = false; int frameIndex = 0; GLRenderToTexture* renderTexture =0; //#include "BroadphaseBenchmark.h" @@ -82,7 +83,7 @@ int selectedDemo = 0; GpuDemo::CreateFunc* allDemos[]= { //ConcaveCompound2Scene::MyCreateFunc, -// GpuConvexScene::MyCreateFunc, + //ConcaveSphereScene::MyCreateFunc, @@ -91,23 +92,24 @@ GpuDemo::CreateFunc* allDemos[]= // ConcaveSphereScene::MyCreateFunc, - + ConcaveScene::MyCreateFunc, GpuBoxPlaneScene::MyCreateFunc, GpuConstraintsDemo::MyCreateFunc, //GpuConvexPlaneScene::MyCreateFunc, + GpuConvexScene::MyCreateFunc, + GpuCompoundScene::MyCreateFunc, GpuSphereScene::MyCreateFunc, - ConcaveSphereScene::MyCreateFunc, + ConcaveScene::MyCreateFunc, - - - + ConcaveSphereScene::MyCreateFunc, + ConcaveCompoundScene::MyCreateFunc, // GpuCompoundPlaneScene::MyCreateFunc, @@ -228,9 +230,15 @@ void MyKeyboardCallback(int key, int state) { window->setRequestExit(); } + if (key==B3G_F2) + { + if (state) + exportMovie = !exportMovie; + } if (key==B3G_F1) { - exportFrame = true; + if (state) + exportFrame = true; } if (sDemo) sDemo->keyboardCallback(key,state); @@ -671,7 +679,7 @@ int main(int argc, char* argv[]) bool useGpu = false; - int maxObjectCapacity=128*1024; + int maxObjectCapacity=512*1024; maxObjectCapacity = b3Max(maxObjectCapacity,ci.arraySizeX*ci.arraySizeX*ci.arraySizeX+10); { @@ -786,7 +794,7 @@ int main(int argc, char* argv[]) assert(err==GL_NO_ERROR); - if (exportFrame) + if (exportFrame || exportMovie) { if (!renderTexture) @@ -808,7 +816,7 @@ int main(int argc, char* argv[]) //glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE); //glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE); - renderTexture->init(g_OpenGLWidth,g_OpenGLHeight,renderTextureId, false); + renderTexture->init(g_OpenGLWidth,g_OpenGLHeight,renderTextureId, RENDERTEXTURE_COLOR); } bool result = renderTexture->enable(); @@ -870,7 +878,7 @@ int main(int argc, char* argv[]) */ - if (exportFrame) + if (exportFrame || exportMovie) { char fileName[1024]; diff --git a/Demos3/GpuDemos/rigidbody/ConcaveScene.cpp b/Demos3/GpuDemos/rigidbody/ConcaveScene.cpp index 6d8bd9cb2..3617ba5eb 100644 --- a/Demos3/GpuDemos/rigidbody/ConcaveScene.cpp +++ b/Demos3/GpuDemos/rigidbody/ConcaveScene.cpp @@ -19,11 +19,11 @@ #include "Bullet3Common/b3Transform.h" #include "Bullet3OpenCL/NarrowphaseCollision/b3ConvexUtility.h" - +#include "../gwenUserInterface.h" #include "OpenGLWindow/GLInstanceGraphicsShape.h" -#define CONCAVE_GAPX 16 +#define CONCAVE_GAPX 14 #define CONCAVE_GAPY 8 -#define CONCAVE_GAPZ 16 +#define CONCAVE_GAPZ 14 GLInstanceGraphicsShape* createGraphicsShapeFromWavefrontObj(std::vector& shapes) @@ -224,7 +224,7 @@ void ConcaveScene::setupScene(const ConstructionInfo& ci) b3Vector3 shift1(0,0,0);//0,230,80);//150,-100,-120); - b3Vector4 scaling(4,4,4,1); + b3Vector4 scaling(10,10,10,1); // createConcaveMesh(ci,"plane100.obj",shift1,scaling); //createConcaveMesh(ci,"plane100.obj",shift,scaling); @@ -269,6 +269,11 @@ void ConcaveScene::setupScene(const ConstructionInfo& ci) m_instancingRenderer->setCameraPitch(45); m_instancingRenderer->setCameraTargetPosition(camPos); m_instancingRenderer->setCameraDistance(155); + char msg[1024]; + int numInstances = m_data->m_rigidBodyPipeline->getNumBodies(); + sprintf(msg,"Num objects = %d",numInstances); + if (ci.m_gui) + ci.m_gui->setStatusBarMessage(msg,true); } diff --git a/Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp b/Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp index 5fd1f59c0..9eb702793 100644 --- a/Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp +++ b/Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp @@ -154,7 +154,7 @@ int GpuConvexScene::createDynamicsObjects2(const ConstructionInfo& ci, const flo { //mass=0.f; } - b3Vector3 position((j&1)+i*2.2,1+j*2.,(j&1)+k*2.2); + b3Vector3 position(((j+1)&1)+i*2.2,1+j*2.,((j+1)&1)+k*2.2); //b3Vector3 position(i*2.2,10+j*1.9,k*2.2); b3Quaternion orn(0,0,0,1); diff --git a/btgui/OpenGLWindow/GLInstancingRenderer.cpp b/btgui/OpenGLWindow/GLInstancingRenderer.cpp index dc071c808..0e64aae4d 100644 --- a/btgui/OpenGLWindow/GLInstancingRenderer.cpp +++ b/btgui/OpenGLWindow/GLInstancingRenderer.cpp @@ -129,7 +129,9 @@ struct InternalDataRenderer : public GLInstanceRendererInternalData float m_mouseYpos; bool m_mouseInitialized; int m_leftMouseButton; + int m_middleMouseButton; int m_rightMouseButton; + int m_altPressed; int m_controlPressed; @@ -150,6 +152,7 @@ struct InternalDataRenderer : public GLInstanceRendererInternalData m_ele(25.f), m_mouseInitialized(false), m_leftMouseButton(0), + m_middleMouseButton(0), m_rightMouseButton(0), m_shadowMap(0), m_shadowTexture(0), @@ -209,6 +212,16 @@ struct InternalDataRenderer : public GLInstanceRendererInternalData // { m_ele += yDelta*MOUSE_MOVE_MULTIPLIER; // } + } + if (m_middleMouseButton) + { + m_cameraTargetPosition += m_cameraUp * yDelta; + + b3Vector3 fwd = m_cameraTargetPosition-m_cameraPosition; + b3Vector3 side = m_cameraUp.cross(fwd); + side.normalize(); + m_cameraTargetPosition += side * xDelta; + } if (m_rightMouseButton) { @@ -234,6 +247,9 @@ struct InternalDataRenderer : public GLInstanceRendererInternalData if (button==0) m_leftMouseButton=state; + if (button==1) + m_middleMouseButton=state; + if (button==2) m_rightMouseButton=state; diff --git a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.cpp b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.cpp index bdf9359c7..3e63fb0b8 100644 --- a/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.cpp +++ b/src/Bullet3OpenCL/BroadphaseCollision/b3GpuSapBroadphase.cpp @@ -54,8 +54,8 @@ m_removedCountGPU(ctx,q) 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); - cl_program sapFastProg = b3OpenCLUtils::compileCLProgramFromString(m_context,m_device,0,&errNum,"",B3_BROADPHASE_SAPFAST_PATH,true); + cl_program sapFastProg = b3OpenCLUtils::compileCLProgramFromString(m_context,m_device,sapFastSrc,&errNum,"",B3_BROADPHASE_SAPFAST_PATH); + //cl_program sapFastProg = b3OpenCLUtils::compileCLProgramFromString(m_context,m_device,0,&errNum,"",B3_BROADPHASE_SAPFAST_PATH,true); b3Assert(errNum==CL_SUCCESS); #ifndef __APPLE__ m_prefixScanFloat4 = new b3PrefixScanFloat4CL(m_context,m_device,m_queue); diff --git a/src/Bullet3OpenCL/BroadphaseCollision/kernels/sapFastKernels.h b/src/Bullet3OpenCL/BroadphaseCollision/kernels/sapFastKernels.h index b9986a74c..24f36bf9a 100644 --- a/src/Bullet3OpenCL/BroadphaseCollision/kernels/sapFastKernels.h +++ b/src/Bullet3OpenCL/BroadphaseCollision/kernels/sapFastKernels.h @@ -32,6 +32,22 @@ static const char* sapFastCL= \ " };\n" "} btAabbCL;\n" "\n" +"typedef struct \n" +"{\n" +" union\n" +" {\n" +" unsigned int m_key;\n" +" unsigned int x;\n" +" };\n" +"\n" +" union\n" +" {\n" +" unsigned int m_value;\n" +" unsigned int y;\n" +" \n" +" };\n" +"}b3SortData;\n" +"\n" "\n" "/// conservative test for overlap between two aabbs\n" "bool TestAabbAgainstAabb2(const btAabbCL* aabb1, __local const btAabbCL* aabb2);\n" @@ -48,6 +64,260 @@ static const char* sapFastCL= \ " return overlap;\n" "}\n" "\n" +"__kernel void computePairsIncremental3dSapKernel( __global const uint2* objectMinMaxIndexGPUaxis0,\n" +" __global const uint2* objectMinMaxIndexGPUaxis1,\n" +" __global const uint2* objectMinMaxIndexGPUaxis2,\n" +" __global const uint2* objectMinMaxIndexGPUaxis0prev,\n" +" __global const uint2* objectMinMaxIndexGPUaxis1prev,\n" +" __global const uint2* objectMinMaxIndexGPUaxis2prev,\n" +" __global const b3SortData* sortedAxisGPU0,\n" +" __global const b3SortData* sortedAxisGPU1,\n" +" __global const b3SortData* sortedAxisGPU2,\n" +" __global const b3SortData* sortedAxisGPU0prev,\n" +" __global const b3SortData* sortedAxisGPU1prev,\n" +" __global const b3SortData* sortedAxisGPU2prev,\n" +" __global int2* addedHostPairsGPU,\n" +" __global int2* removedHostPairsGPU,\n" +" volatile __global int* addedHostPairsCount,\n" +" volatile __global int* removedHostPairsCount,\n" +" int maxCapacity,\n" +" int numObjects)\n" +"{\n" +" int i = get_global_id(0);\n" +" if (i>=numObjects)\n" +" return;\n" +"\n" +" __global const uint2* objectMinMaxIndexGPU[3][2];\n" +" objectMinMaxIndexGPU[0][0]=objectMinMaxIndexGPUaxis0;\n" +" objectMinMaxIndexGPU[1][0]=objectMinMaxIndexGPUaxis1;\n" +" objectMinMaxIndexGPU[2][0]=objectMinMaxIndexGPUaxis2;\n" +" objectMinMaxIndexGPU[0][1]=objectMinMaxIndexGPUaxis0prev;\n" +" objectMinMaxIndexGPU[1][1]=objectMinMaxIndexGPUaxis1prev;\n" +" objectMinMaxIndexGPU[2][1]=objectMinMaxIndexGPUaxis2prev;\n" +"\n" +" __global const b3SortData* sortedAxisGPU[3][2];\n" +" sortedAxisGPU[0][0] = sortedAxisGPU0;\n" +" sortedAxisGPU[1][0] = sortedAxisGPU1;\n" +" sortedAxisGPU[2][0] = sortedAxisGPU2;\n" +" sortedAxisGPU[0][1] = sortedAxisGPU0prev;\n" +" sortedAxisGPU[1][1] = sortedAxisGPU1prev;\n" +" sortedAxisGPU[2][1] = sortedAxisGPU2prev;\n" +"\n" +" int m_currentBuffer = 0;\n" +"\n" +" for (int axis=0;axis<3;axis++)\n" +" {\n" +" //int i = checkObjects[a];\n" +"\n" +" unsigned int curMinIndex = objectMinMaxIndexGPU[axis][m_currentBuffer][i].x;\n" +" unsigned int curMaxIndex = objectMinMaxIndexGPU[axis][m_currentBuffer][i].y;\n" +" unsigned int prevMinIndex = objectMinMaxIndexGPU[axis][1-m_currentBuffer][i].x;\n" +" int dmin = curMinIndex - prevMinIndex;\n" +" \n" +" unsigned int prevMaxIndex = objectMinMaxIndexGPU[axis][1-m_currentBuffer][i].y;\n" +"\n" +" int dmax = curMaxIndex - prevMaxIndex;\n" +" \n" +" for (int otherbuffer = 0;otherbuffer<2;otherbuffer++)\n" +" {\n" +" if (dmin!=0)\n" +" {\n" +" int stepMin = dmin<0 ? -1 : 1;\n" +" for (int j=prevMinIndex;j!=curMinIndex;j+=stepMin)\n" +" {\n" +" int otherIndex2 = sortedAxisGPU[axis][otherbuffer][j].y;\n" +" int otherIndex = otherIndex2/2;\n" +" if (otherIndex!=i)\n" +" {\n" +" bool otherIsMax = ((otherIndex2&1)!=0);\n" +"\n" +" if (otherIsMax)\n" +" {\n" +" \n" +" bool overlap = true;\n" +"\n" +" for (int ax=0;ax<3;ax++)\n" +" {\n" +" if ((objectMinMaxIndexGPU[ax][m_currentBuffer][i].x > objectMinMaxIndexGPU[ax][m_currentBuffer][otherIndex].y) ||\n" +" (objectMinMaxIndexGPU[ax][m_currentBuffer][i].y < objectMinMaxIndexGPU[ax][m_currentBuffer][otherIndex].x))\n" +" overlap=false;\n" +" }\n" +"\n" +" // b3Assert(overlap2==overlap);\n" +"\n" +" bool prevOverlap = true;\n" +"\n" +" for (int ax=0;ax<3;ax++)\n" +" {\n" +" if ((objectMinMaxIndexGPU[ax][1-m_currentBuffer][i].x > objectMinMaxIndexGPU[ax][1-m_currentBuffer][otherIndex].y) ||\n" +" (objectMinMaxIndexGPU[ax][1-m_currentBuffer][i].y < objectMinMaxIndexGPU[ax][1-m_currentBuffer][otherIndex].x))\n" +" prevOverlap=false;\n" +" }\n" +" \n" +"\n" +" //b3Assert(overlap==overlap2);\n" +" \n" +"\n" +"\n" +" if (dmin<0)\n" +" {\n" +" if (overlap && !prevOverlap)\n" +" {\n" +" //add a pair\n" +" int2 newPair;\n" +" if (i<=otherIndex)\n" +" {\n" +" newPair.x = i;\n" +" newPair.y = otherIndex;\n" +" } else\n" +" {\n" +" newPair.x = otherIndex;\n" +" newPair.y = i;\n" +" }\n" +" \n" +" {\n" +" int curPair = atomic_inc(addedHostPairsCount);\n" +" if (curPair objectMinMaxIndexGPU[ax][m_currentBuffer][otherIndex].y) ||\n" +" (objectMinMaxIndexGPU[ax][m_currentBuffer][i].y < objectMinMaxIndexGPU[ax][m_currentBuffer][otherIndex].x))\n" +" overlap=false;\n" +" }\n" +" //b3Assert(overlap2==overlap);\n" +"\n" +" bool prevOverlap = true;\n" +"\n" +" for (int ax=0;ax<3;ax++)\n" +" {\n" +" if ((objectMinMaxIndexGPU[ax][1-m_currentBuffer][i].x > objectMinMaxIndexGPU[ax][1-m_currentBuffer][otherIndex].y) ||\n" +" (objectMinMaxIndexGPU[ax][1-m_currentBuffer][i].y < objectMinMaxIndexGPU[ax][1-m_currentBuffer][otherIndex].x))\n" +" prevOverlap=false;\n" +" }\n" +" \n" +"\n" +" if (dmax>0)\n" +" {\n" +" if (overlap && !prevOverlap)\n" +" {\n" +" //add a pair\n" +" int2 newPair;\n" +" if (i<=otherIndex)\n" +" {\n" +" newPair.x = i;\n" +" newPair.y = otherIndex;\n" +" } else\n" +" {\n" +" newPair.x = otherIndex;\n" +" newPair.y = i;\n" +" }\n" +" {\n" +" int curPair = atomic_inc(addedHostPairsCount);\n" +" if (curPair min?\n" +" //remove a pair\n" +" int2 removedPair;\n" +" if (i<=otherIndex)\n" +" {\n" +" removedPair.x = i;\n" +" removedPair.y = otherIndex;\n" +" } else\n" +" {\n" +" removedPair.x = otherIndex;\n" +" removedPair.y = i;\n" +" }\n" +" {\n" +" int curPair = atomic_inc(removedHostPairsCount);\n" +" if (curPair