diff --git a/demo/gpudemo/GpuDemo.cpp b/demo/gpudemo/GpuDemo.cpp index 32805bedc..f9c2a21c7 100644 --- a/demo/gpudemo/GpuDemo.cpp +++ b/demo/gpudemo/GpuDemo.cpp @@ -40,9 +40,9 @@ void GpuDemo::initCL(int preferredDeviceIndex, int preferredPlatformIndex) int ciErrNum = 0; //#ifdef CL_PLATFORM_INTEL - cl_device_type deviceType = CL_DEVICE_TYPE_ALL; + //cl_device_type deviceType = CL_DEVICE_TYPE_ALL; //#else - //cl_device_type deviceType = CL_DEVICE_TYPE_CPU; + cl_device_type deviceType = CL_DEVICE_TYPE_GPU; //#endif cl_platform_id platformId; diff --git a/demo/gpudemo/GpuDemo.h b/demo/gpudemo/GpuDemo.h index 5c30cf015..c321971aa 100644 --- a/demo/gpudemo/GpuDemo.h +++ b/demo/gpudemo/GpuDemo.h @@ -37,13 +37,13 @@ public: :useOpenCL(true), preferredOpenCLPlatformIndex(-1), preferredOpenCLDeviceIndex(-1), - arraySizeX(10), - arraySizeY(10 ), - arraySizeZ(10), + arraySizeX(20), + arraySizeY(20 ), + arraySizeZ(20), m_useConcaveMesh(false), - gapX(6.3), - gapY(12.0), - gapZ(6.3), + gapX(4.3), + gapY(4.0), + gapZ(4.3), m_instancingRenderer(0), m_window(0) { diff --git a/demo/gpudemo/main_opengl3core.cpp b/demo/gpudemo/main_opengl3core.cpp index bfee4824e..04c0f0e77 100644 --- a/demo/gpudemo/main_opengl3core.cpp +++ b/demo/gpudemo/main_opengl3core.cpp @@ -552,7 +552,7 @@ int main(int argc, char* argv[]) bool useGpu = false; - int maxObjectCapacity=128*1024; + int maxObjectCapacity=256*1024; ci.m_instancingRenderer = new GLInstancingRenderer(maxObjectCapacity);//render.getInstancingRenderer(); ci.m_window = window; diff --git a/demo/gpudemo/rigidbody/GpuCompoundScene.cpp b/demo/gpudemo/rigidbody/GpuCompoundScene.cpp index f0e62e0e5..ed1d71379 100644 --- a/demo/gpudemo/rigidbody/GpuCompoundScene.cpp +++ b/demo/gpudemo/rigidbody/GpuCompoundScene.cpp @@ -21,74 +21,6 @@ #include "OpenGLWindow/GLInstanceGraphicsShape.h" -#if 0 -GraphicsShape* createGraphicsShapeFromCompoundShape(btCompoundShape* compound) -{ - GraphicsShape* gfxShape = new GraphicsShape(); - btAlignedObjectArray* vertexArray = new btAlignedObjectArray; - btAlignedObjectArray* indexArray = new btAlignedObjectArray; - - - - //create a graphics shape for each child, combine them into a single graphics shape using their child transforms - for (int i=0;igetNumChildShapes();i++) - { - btAssert(compound->getChildShape(i)->isPolyhedral()); - if (compound->getChildShape(i)->isPolyhedral()) - { - btPolyhedralConvexShape* convexHull = (btPolyhedralConvexShape*) compound->getChildShape(i); - btTransform tr = compound->getChildTransform(i); - - const btConvexPolyhedron* polyhedron = convexHull->getConvexPolyhedron(); - GraphicsShape* childGfxShape = createGraphicsShapeFromConvexHull(polyhedron); - int baseIndex = vertexArray->size(); - - for (int j=0;jm_numIndices;j++) - indexArray->push_back(childGfxShape->m_indices[j]+baseIndex); - - GLInstanceVertex* orgVerts = (GLInstanceVertex*)childGfxShape->m_vertices; - - for (int j=0;jm_numvertices;j++) - { - GLInstanceVertex vtx; - btVector3 pos(orgVerts[j].xyzw[0],orgVerts[j].xyzw[1],orgVerts[j].xyzw[2]); - pos = tr*pos; - vtx.xyzw[0] = childGfxShape->m_scaling[0]*pos.x(); - vtx.xyzw[1] = childGfxShape->m_scaling[1]*pos.y(); - vtx.xyzw[2] = childGfxShape->m_scaling[2]*pos.z(); - vtx.xyzw[3] = 10.f; - - vtx.uv[0] = 0.5f; - vtx.uv[1] = 0.5f; - - btVector3 normal(orgVerts[j].normal[0],orgVerts[j].normal[1],orgVerts[j].normal[2]); - normal = tr.getBasis()*normal; - vtx.normal[0] = normal.x(); - vtx.normal[1] = normal.y(); - vtx.normal[2] = normal.z(); - vertexArray->push_back(vtx); - } - } - } - - btPolyhedralConvexShape* convexHull = (btPolyhedralConvexShape*) compound->getChildShape(0); - const btConvexPolyhedron* polyhedron = convexHull->getConvexPolyhedron(); - GraphicsShape* childGfxShape = createGraphicsShapeFromConvexHull(polyhedron); - - gfxShape->m_indices = &indexArray->at(0); - gfxShape->m_numIndices = indexArray->size(); - gfxShape->m_vertices = &vertexArray->at(0).xyzw[0]; - gfxShape->m_numvertices = vertexArray->size(); - gfxShape->m_scaling[0] = 1; - gfxShape->m_scaling[1] = 1; - gfxShape->m_scaling[2] = 1; - gfxShape->m_scaling[3] = 1; - - return gfxShape; -} -#endif - - void GpuCompoundScene::setupScene(const ConstructionInfo& ci) { diff --git a/demo/gpudemo/rigidbody/GpuConvexScene.cpp b/demo/gpudemo/rigidbody/GpuConvexScene.cpp index df9e69174..454605d42 100644 --- a/demo/gpudemo/rigidbody/GpuConvexScene.cpp +++ b/demo/gpudemo/rigidbody/GpuConvexScene.cpp @@ -27,35 +27,62 @@ void GpuConvexScene::setupScene(const ConstructionInfo& ci) int group=1; int mask=1; int index=10; - float scaling[4] = {1,1,1,1}; + - int colIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling); - - - for (int i=0;im_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling); + btVector3 position(0,0,0); + btQuaternion orn(1,0,0,0); + + btVector4 color(0,0,1,1); + + int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling); + int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(0.f,position,orn,colIndex,index); + + index++; + } - btVector3 position(i*ci.gapX,j*ci.gapY,k*ci.gapZ); - btQuaternion orn(1,0,0,0); + + + { + btVector4 colors[4] = + { + btVector4(1,0,0,1), + btVector4(0,1,0,1), + btVector4(0,1,1,1), + btVector4(1,1,0,1), + }; + + int curColor = 0; + float scaling[4] = {1,1,1,1}; + int colIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling); + for (int i=0;iregisterGraphicsInstance(shapeId,position,orn,color,scaling); - int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(mass,position,orn,colIndex,index); + btVector4 color = colors[curColor]; + curColor++; + curColor&=3; + btVector4 scaling(1,1,1,1); + int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling); + int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(mass,position,orn,colIndex,index); - index++; + index++; + } } } } - - float camPos[4]={65.5,4.5,65.5,0}; + float camPos[4]={ci.arraySizeX,ci.arraySizeY/2,ci.arraySizeZ,0}; //float camPos[4]={1,12.5,1.5,0}; m_instancingRenderer->setCameraTargetPosition(camPos); - m_instancingRenderer->setCameraDistance(90); + m_instancingRenderer->setCameraDistance(120); } \ No newline at end of file diff --git a/opencl/gpu_rigidbody/host/btConfig.h b/opencl/gpu_rigidbody/host/btConfig.h index 3f44e410b..32a6cbc5c 100644 --- a/opencl/gpu_rigidbody/host/btConfig.h +++ b/opencl/gpu_rigidbody/host/btConfig.h @@ -18,7 +18,7 @@ struct btConfig int m_maxTriConvexPairCapacity; btConfig() - :m_maxConvexBodies(16*1024), + :m_maxConvexBodies(32*1024), m_maxConvexShapes(8192), m_maxVerticesPerFace(64), m_maxFacesPerShape(64), @@ -26,7 +26,7 @@ struct btConfig m_maxConvexIndices(8192), m_maxConvexUniqueEdges(8192), m_maxCompoundChildShapes(8192), - m_maxTriConvexPairCapacity(16*1024) + m_maxTriConvexPairCapacity(64*1024) { m_maxBroadphasePairs = 16*m_maxConvexBodies; } diff --git a/opencl/gpu_rigidbody/host/btGpuBatchingPgsSolver.cpp b/opencl/gpu_rigidbody/host/btGpuBatchingPgsSolver.cpp index d60fc1a23..03b56d9fe 100644 --- a/opencl/gpu_rigidbody/host/btGpuBatchingPgsSolver.cpp +++ b/opencl/gpu_rigidbody/host/btGpuBatchingPgsSolver.cpp @@ -35,7 +35,7 @@ enum }; -bool gpuBatchContacts = false; +bool gpuBatchContacts = true; bool gpuSolveConstraint = true; diff --git a/opencl/gpu_rigidbody/host/btGpuRigidBodyPipeline.cpp b/opencl/gpu_rigidbody/host/btGpuRigidBodyPipeline.cpp index cd2ada66e..ab2e6a15a 100644 --- a/opencl/gpu_rigidbody/host/btGpuRigidBodyPipeline.cpp +++ b/opencl/gpu_rigidbody/host/btGpuRigidBodyPipeline.cpp @@ -143,10 +143,13 @@ void btGpuRigidBodyPipeline::integrate(float timeStep) launcher.setBuffer(m_data->m_narrowphase->getBodiesGpu()); int numBodies = m_data->m_narrowphase->getNumBodiesGpu(); launcher.setConst(numBodies); - launcher.setConst(timeStep); float angularDamp = 0.99f; launcher.setConst(angularDamp); + + btVector3 gravity(0.f,-9.8f,0.f); + launcher.setConst(gravity); + launcher.launch1D(numBodies); } diff --git a/opencl/gpu_rigidbody/kernels/integrateKernel.cl b/opencl/gpu_rigidbody/kernels/integrateKernel.cl index c96f5b696..96f5e4d64 100644 --- a/opencl/gpu_rigidbody/kernels/integrateKernel.cl +++ b/opencl/gpu_rigidbody/kernels/integrateKernel.cl @@ -42,7 +42,7 @@ typedef struct __kernel void - integrateTransformsKernel( __global Body* bodies,const int numNodes, float timeStep, float angularDamping) + integrateTransformsKernel( __global Body* bodies,const int numNodes, float timeStep, float angularDamping, float4 gravityAcceleration) { int nodeID = get_global_id(0); float BT_GPU_ANGULAR_MOTION_THRESHOLD = (0.25f * 3.14159254f); @@ -86,7 +86,6 @@ __kernel void bodies[nodeID].m_pos += bodies[nodeID].m_linVel * timeStep; //apply gravity - float4 gravityAcceleration = (float4)(0.f,-9.8f,0.f,0.f); bodies[nodeID].m_linVel += gravityAcceleration * timeStep; } diff --git a/opencl/gpu_rigidbody/kernels/integrateKernel.h b/opencl/gpu_rigidbody/kernels/integrateKernel.h index 256f0697d..296234c81 100644 --- a/opencl/gpu_rigidbody/kernels/integrateKernel.h +++ b/opencl/gpu_rigidbody/kernels/integrateKernel.h @@ -44,7 +44,7 @@ static const char* integrateKernelCL= \ "\n" "\n" "__kernel void \n" -" integrateTransformsKernel( __global Body* bodies,const int numNodes, float timeStep, float angularDamping)\n" +" integrateTransformsKernel( __global Body* bodies,const int numNodes, float timeStep, float angularDamping, float4 gravityAcceleration)\n" "{\n" " int nodeID = get_global_id(0);\n" " float BT_GPU_ANGULAR_MOTION_THRESHOLD = (0.25f * 3.14159254f);\n" @@ -88,7 +88,6 @@ static const char* integrateKernelCL= \ " bodies[nodeID].m_pos += bodies[nodeID].m_linVel * timeStep;\n" " \n" " //apply gravity\n" -" float4 gravityAcceleration = (float4)(0.f,-9.8f,0.f,0.f);\n" " bodies[nodeID].m_linVel += gravityAcceleration * timeStep;\n" " \n" " }\n"