diff --git a/Demos/Gpu3dDemo/BasicDemo3d.cpp b/Demos/Gpu3dDemo/BasicDemo3d.cpp new file mode 100644 index 000000000..52dcb37e4 --- /dev/null +++ b/Demos/Gpu3dDemo/BasicDemo3d.cpp @@ -0,0 +1,844 @@ +/* +Bullet Continuous Collision Detection and Physics Library +Copyright (c) 2003-2006 Erwin Coumans http://continuousphysics.com/Bullet/ + +This software is provided 'as-is', without any express or implied warranty. +In no event will the authors be held liable for any damages arising from the use of this software. +Permission is granted to anyone to use this software for any purpose, +including commercial applications, and to alter it and redistribute it freely, +subject to the following restrictions: + +1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. +2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. +3. This notice may not be removed or altered from any source distribution. +*/ + +#include "BulletMultiThreaded/SpuNarrowPhaseCollisionTask/SpuGatheringCollisionTask.h" +#include "BulletMultiThreaded/SpuContactManifoldCollisionAlgorithm.h" + +#include "btGpuDemoDynamicsWorld3D.h" + +#include "BulletMultiThreaded/SpuGatheringCollisionDispatcher.h" +#include "BulletMultiThreaded/Win32ThreadSupport.h" +#include "GLDebugFont.h" +//@ extern int gSkippedCol; +//@ extern int gProcessedCol; + + + +#define SPEC_TEST 0 + +#ifdef _DEBUG + #define LARGE_DEMO 0 +// #define LARGE_DEMO 1 +#else + #define LARGE_DEMO 1 +#endif + +#if LARGE_DEMO + ///create 512 (8x8x8) dynamic object +// #define ARRAY_SIZE_X 100 +// #define ARRAY_SIZE_Y 100 +// #define ARRAY_SIZE_Z 1 +// #define ARRAY_SIZE_X 228 +// #define ARRAY_SIZE_Y 228 +// #define ARRAY_SIZE_X 30 +// #define ARRAY_SIZE_Y 100 + +#define ARRAY_SIZE_X 8 +#define ARRAY_SIZE_Y 47 +#define ARRAY_SIZE_Z 8 +#else + ///create 125 (5x5x5) dynamic object + #define ARRAY_SIZE_X 45 + #define ARRAY_SIZE_Y 45 +// #define ARRAY_SIZE_Z 5 + #define ARRAY_SIZE_Z 1 +#endif + + +//maximum number of objects (and allow user to shoot additional boxes) +#define NUM_SMALL_PROXIES (ARRAY_SIZE_X*ARRAY_SIZE_Y*ARRAY_SIZE_Z) +#define MAX_PROXIES (NUM_SMALL_PROXIES + 1024) +#define MAX_LARGE_PROXIES 10 +#define MAX_SMALL_PROXIES (MAX_PROXIES - MAX_LARGE_PROXIES) + +///scaling of the objects (0.1 = 20 centimeter boxes ) +//#define SCALING 0.1 +#define SCALING 1 +#define START_POS_X 0 +#define START_POS_Y 5 +#define START_POS_Z 0 + +#include "BasicDemo3D.h" +#include "GlutStuff.h" +///btBulletDynamicsCommon.h is the main Bullet include file, contains most common include files. +#include "btBulletDynamicsCommon.h" +#include //printf debugging + +#include "BulletCollision/CollisionDispatch/btSimulationIslandManager.h" + +#include "../Extras/CUDA/btCudaBroadphase.h" + +btScalar gTimeStep = btScalar(1./60.); + +bool gbDrawBatches = false; +int gSelectedBatch = CUDA_DEMO_DYNAMICS_WORLD3D_MAX_BATCHES; +bool gUseCPUSolver = false; +bool gUseSolver2 = true; +bool gDrawWire = false; +bool gUseCudaMotIntegr = true; + + +void BasicDemo3D::clientMoveAndDisplay() +{ + glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); + + //simple dynamics world doesn't handle fixed-time-stepping + float ms = getDeltaTimeMicroseconds(); + + ///step the simulation + if (m_dynamicsWorld) + { +// btCudaDemoPairCache* pc = (btCudaDemoPairCache*)m_dynamicsWorld->getPairCache(); +// pc->m_numSmallProxies = m_dynamicsWorld->getNumCollisionObjects(); // - 1; // exclude floor + m_dynamicsWorld->stepSimulation(gTimeStep,0);//ms / 1000000.f); + //optional but useful: debug drawing + m_dynamicsWorld->debugDrawWorld(); + } + renderme(); + + ms = getDeltaTimeMicroseconds(); + + glFlush(); + + glutSwapBuffers(); + +} + + + +void BasicDemo3D::displayCallback(void) { + + glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); + + renderme(); + + //optional but useful: debug drawing to detect problems + if (m_dynamicsWorld) + m_dynamicsWorld->debugDrawWorld(); + + glFlush(); + glutSwapBuffers(); +} + +#define NUM_SOLVERS 11 +static btConstraintSolver* sConstraintSolvers[NUM_SOLVERS]; +static int sCurrSolverIndex = 9; +static char* sConstraintSolverNames[NUM_SOLVERS] = +{ + "btSequentialImpulseConstraintSolver", + "btParallelBatchConstraintSolver", + "btCudaConstraintSolver", + "btParallelBatchConstraintSolver2", + "btParallelBatchConstraintSolver3", + "btCudaConstraintSolver3", + "btParallelBatchConstraintSolver4", + "btCudaConstraintSolver4", + "btParallelBatchConstraintSolver5", + "btParallelBatchConstraintSolver6", + "btCudaConstraintSolver6", +}; + +//btVector3 gWorldMin(-228,-228,-32); +//btVector3 gWorldMin(-228,0,-32); +//btVector3 gWorldMax(228,228,32); + +//btVector3 gWorldMin(-150,-228,-32); +//btVector3 gWorldMax(150,228,32); + +#define POS_OFFS_X (ARRAY_SIZE_X * SCALING + 50) +#define POS_OFFS_Y (ARRAY_SIZE_Y * SCALING ) +#define POS_OFFS_Z (ARRAY_SIZE_Z * SCALING + 5) + +btVector3 gWorldMin(-POS_OFFS_X, -ARRAY_SIZE_Y*SCALING, -80-POS_OFFS_Z); +btVector3 gWorldMax( POS_OFFS_X, POS_OFFS_Y, 80+POS_OFFS_Z); + +//btCudaDemoPairCache* gPairCache; +btHashedOverlappingPairCache* gPairCache; + +void BasicDemo3D::initPhysics() +{ + setTexturing(true); + setShadows(false); + +// setCameraDistance(btScalar(SCALING*50.)); +#if LARGE_DEMO + setCameraDistance(btScalar(SCALING*50.)); +#else + setCameraDistance(btScalar(SCALING*20.)); +#endif + + m_cameraTargetPosition.setValue(START_POS_X, -START_POS_Y-20, START_POS_Z); + m_azi = btScalar(0.f); + m_ele = btScalar(0.f); + + ///collision configuration contains default setup for memory, collision setup + + btDefaultCollisionConstructionInfo dci; + dci.m_defaultMaxPersistentManifoldPoolSize=100000; + dci.m_defaultMaxCollisionAlgorithmPoolSize=100000; + + ///SpuContactManifoldCollisionAlgorithm is larger than any of the other collision algorithms +//@@ dci.m_customMaxCollisionAlgorithmSize = sizeof(SpuContactManifoldCollisionAlgorithm); + + m_collisionConfiguration = new btDefaultCollisionConfiguration(dci); + + ///use the default collision dispatcher. For parallel processing you can use a diffent dispatcher (see Extras/BulletMultiThreaded) + //m_dispatcher = new btCollisionDispatcher(m_collisionConfiguration); +#ifdef SINGLE_THREADED_NARROWPHASE + m_dispatcher = new btCollisionDispatcher(m_collisionConfiguration); +#else + unsigned int maxNumOutstandingTasks =4; + //createCollisionLocalStoreMemory(); + //processSolverTask + Win32ThreadSupport::Win32ThreadConstructionInfo threadConstructionInfo("narrowphase_multi",processCollisionTask,createCollisionLocalStoreMemory,maxNumOutstandingTasks); + class btThreadSupportInterface* threadInterface = new Win32ThreadSupport(threadConstructionInfo); + m_dispatcher = new SpuGatheringCollisionDispatcher(threadInterface,maxNumOutstandingTasks,m_collisionConfiguration); +#endif //SINGLE_THREADED_NARROWPHASE + + +//## m_dispatcher->registerCollisionCreateFunc(BOX_SHAPE_PROXYTYPE,BOX_SHAPE_PROXYTYPE,new btEmptyAlgorithm::CreateFunc); +//## m_dispatcher->registerCollisionCreateFunc(CUSTOM_CONVEX_SHAPE_TYPE,CUSTOM_CONVEX_SHAPE_TYPE,new btBox2dBox2dCollisionAlgorithm::CreateFunc); + +// m_broadphase = new btDbvtBroadphase(); + + +//## gPairCache = new (btAlignedAlloc(sizeof(btCudaDemoPairCache),16)) btCudaDemoPairCache(MAX_PROXIES, 24, MAX_SMALL_PROXIES); +// gPairCache = NULL; + gPairCache = new (btAlignedAlloc(sizeof(btHashedOverlappingPairCache),16)) btHashedOverlappingPairCache(); + + //m_broadphase = new btSimpleBroadphase(16384, gPairCache); + +/* +btCudaBroadphase::btCudaBroadphase( btOverlappingPairCache* overlappingPairCache, + const btVector3& worldAabbMin,const btVector3& worldAabbMax, + int gridSizeX, int gridSizeY, int gridSizeZ, + int maxSmallProxies, int maxLargeProxies, int maxPairsPerBody, + int maxBodiesPerCell, + btScalar cellFactorAABB) +*/ +// btVector3 numOfCells = (gWorldMax - gWorldMin) / (2. * SCALING * 0.7); + btVector3 numOfCells = (gWorldMax - gWorldMin) / (2. * SCALING); + int numOfCellsX = (int)numOfCells[0]; + int numOfCellsY = (int)numOfCells[1]; + int numOfCellsZ = (int)numOfCells[2]; + +// m_broadphase = new bt3DGridBroadphase(gPairCache, gWorldMin, gWorldMax,numOfCellsX, numOfCellsY, numOfCellsZ,MAX_SMALL_PROXIES,10,8,8,1./1.5); +//#define USE_CUDA_BROADPHASE 1 +#ifdef USE_CUDA_BROADPHASE + m_broadphase = new btCudaBroadphase(gPairCache, gWorldMin, gWorldMax,numOfCellsX, numOfCellsY, numOfCellsZ,MAX_SMALL_PROXIES,20,18,8,1./1.5); +#else + btDbvtBroadphase* dbvt = new btDbvtBroadphase(gPairCache); + m_broadphase = dbvt; + dbvt->m_deferedcollide=false; + dbvt->m_prediction = 0.f; + +#endif + //m_broadphase = new btAxisSweep3(gWorldMin,gWorldMax,32000,0,true);//(btDbvtBroadphase(gPairCache); + + // create solvers for tests + ///the default constraint solver + sConstraintSolvers[0] = new btSequentialImpulseConstraintSolver(); +/* + sConstraintSolvers[1] = new btParallelBatchConstraintSolver(); + sConstraintSolvers[2] = new btCudaConstraintSolver(); + sConstraintSolvers[3] = new btParallelBatchConstraintSolver2(); + sConstraintSolvers[4] = new btParallelBatchConstraintSolver3(); + sConstraintSolvers[5] = new btCudaConstraintSolver3(); + sConstraintSolvers[6] = new btParallelBatchConstraintSolver4(); + sConstraintSolvers[7] = new btCudaConstraintSolver4(); + sConstraintSolvers[8] = new btParallelBatchConstraintSolver5(); + sConstraintSolvers[9] = new btParallelBatchConstraintSolver6(); + sConstraintSolvers[10] = new btCudaConstraintSolver6(); +*/ + sCurrSolverIndex = 0; + m_solver = sConstraintSolvers[sCurrSolverIndex]; + printf("\nUsing %s\n", sConstraintSolverNames[sCurrSolverIndex]); + +// sCudaMotionInterface = new btCudaMotionInterface(MAX_PROXIES); +// m_dynamicsWorld = new btDiscreteDynamicsWorld(m_dispatcher,m_broadphase,m_solver,m_collisionConfiguration, sCudaMotionInterface); +// m_dynamicsWorld = new btDiscreteDynamicsWorld(m_dispatcher,m_broadphase,m_solver,m_collisionConfiguration); +//## btCudaDemoDynamicsWorld* pDdw = new btCudaDemoDynamicsWorld(m_dispatcher,m_broadphase,m_solver,m_collisionConfiguration); + btCudaDemoDynamicsWorld3D* pDdw = new btCudaDemoDynamicsWorld3D(m_dispatcher,m_broadphase,m_solver,m_collisionConfiguration); + m_dynamicsWorld = pDdw; + pDdw->getDispatchInfo().m_enableSPU=true; + pDdw->getSimulationIslandManager()->setSplitIslands(sCurrSolverIndex == 0); + pDdw->setObjRad(SCALING); + pDdw->setWorldMin(gWorldMin); + pDdw->setWorldMax(gWorldMax); +#ifdef BT_USE_CUDA + gUseCPUSolver = false; +#else + gUseCPUSolver = true; +#endif + pDdw->setUseCPUSolver(gUseCPUSolver); +// pDdw->setUseSolver2(gUseSolver2); + +// m_dynamicsWorld->setGravity(btVector3(0,0,0)); + m_dynamicsWorld->setGravity(btVector3(0.f,-10.f,0.f)); + m_dynamicsWorld->getSolverInfo().m_numIterations = 4; + + + { + //create a few dynamic rigidbodies + // Re-using the same collision is better for memory usage and performance + + + //btCollisionShape* colShape = new btBoxShape(btVector3(SCALING*1,SCALING*1,0.1));//SCALING*1)); +//## btCollisionShape* colShape = new btBox2dShape(btVector3(SCALING*.7,SCALING*.7,0.1));//SCALING*1)); + btCollisionShape* colShape = new btBoxShape(btVector3(SCALING*.7,SCALING*.7, SCALING*.7)); + + //btCollisionShape* colShape = new btSphereShape(btScalar(1.)); + m_collisionShapes.push_back(colShape); + + /// Create Dynamic Objects + btTransform startTransform; + startTransform.setIdentity(); + + btScalar mass(1.f); + + //rigidbody is dynamic if and only if mass is non zero, otherwise static + bool isDynamic = (mass != 0.f); + + btVector3 localInertia(0,0,0); + if (isDynamic) + colShape->calculateLocalInertia(mass,localInertia); +#if (!SPEC_TEST) + float start_x = START_POS_X - ARRAY_SIZE_X * SCALING; + float start_y = START_POS_Y - ARRAY_SIZE_Y * SCALING; + float start_z = START_POS_Z - ARRAY_SIZE_Z * SCALING; + + for (int k=0;kaddRigidBody(body); + } + } + } +#else + // narrowphase test - 2 bodies at the same position + float start_x = START_POS_X; + float start_y = START_POS_Y; + float start_z = START_POS_Z; +// startTransform.setOrigin(SCALING*btVector3(start_x,start_y-14.f,start_z)); + startTransform.setOrigin(SCALING*btVector3(start_x,start_y-11.f,start_z)); + btRigidBody::btRigidBodyConstructionInfo rbInfo(mass,0,colShape,localInertia); + rbInfo.m_startWorldTransform=startTransform; + btRigidBody* body = new btRigidBody(rbInfo); + m_dynamicsWorld->addRigidBody(body); +// startTransform.setOrigin(SCALING*btVector3(start_x+1.2f,start_y+1.4f-14.f,start_z)); + startTransform.setOrigin(SCALING*btVector3(start_x,start_y + 1.5f -11.f, start_z)); + rbInfo.m_startWorldTransform=startTransform; + body = new btRigidBody(rbInfo); + m_dynamicsWorld->addRigidBody(body); +#endif + } + + +#if 0 + ///create a few basic rigid bodies +// btCollisionShape* groundShape = new btBox2dShape(btVector3(btScalar(50.),btScalar(1.),btScalar(50.))); +// btCollisionShape* groundShape = new btBox2dShape(btVector3(btScalar(228.),btScalar(1.),btScalar(228.))); +// btCollisionShape* groundShape = new btBoxShape(btVector3(btScalar(228.),btScalar(1.),btScalar(228.))); +// btCollisionShape* groundShape = new btBoxShape(btVector3(btScalar(50.),btScalar(50.),btScalar(50.))); +// btCollisionShape* groundShape = new btStaticPlaneShape(btVector3(0,1,0),50); + btCollisionShape* groundShape = new btBoxShape(btVector3(POS_OFFS_X, btScalar(1.), POS_OFFS_Z)); + + m_collisionShapes.push_back(groundShape); + + btTransform groundTransform; + groundTransform.setIdentity(); + groundTransform.setOrigin(btVector3(0, gWorldMin[1], 0)); + +// groundTransform.setOrigin(btVector3(0,-5,0)); +// groundTransform.setOrigin(btVector3(0,-50,0)); + + //We can also use DemoApplication::localCreateRigidBody, but for clarity it is provided here: + { + btScalar mass(0.); + + //rigidbody is dynamic if and only if mass is non zero, otherwise static + bool isDynamic = (mass != 0.f); + + btVector3 localInertia(0,0,0); + if (isDynamic) + groundShape->calculateLocalInertia(mass,localInertia); + + //using motionstate is recommended, it provides interpolation capabilities, and only synchronizes 'active' objects + btDefaultMotionState* myMotionState = new btDefaultMotionState(groundTransform); + btRigidBody::btRigidBodyConstructionInfo rbInfo(mass,myMotionState,groundShape,localInertia); + btRigidBody* body = new btRigidBody(rbInfo); + + //add the body to the dynamics world + m_dynamicsWorld->addRigidBody(body); + } +#endif + //clientResetScene(); +} + +void BasicDemo3D::clientResetScene() +{ + DemoApplication::clientResetScene(); + btCudaDemoDynamicsWorld3D* pDdw = (btCudaDemoDynamicsWorld3D*)m_dynamicsWorld; + pDdw->resetScene(); +#if SPEC_TEST + { + float start_x = START_POS_X; + float start_y = START_POS_Y; + float start_z = START_POS_Z; + int numObjects = m_dynamicsWorld->getNumCollisionObjects(); + btCollisionObjectArray& collisionObjects = m_dynamicsWorld->getCollisionObjectArray(); + btTransform startTransform; + startTransform.setIdentity(); + for(int n = 0; n < numObjects; n++) + { + btCollisionObject* colObj = collisionObjects[n]; + btRigidBody* rb = btRigidBody::upcast(colObj); + if(!n) + { +// startTransform.setOrigin(SCALING*btVector3(start_x,start_y-14.f,start_z)); + startTransform.setOrigin(SCALING*btVector3(start_x,start_y-11.f,start_z)); + } + else + { +// startTransform.setOrigin(SCALING*btVector3(start_x+1.2f,start_y+1.4f-14.f,start_z)); + startTransform.setOrigin(SCALING*btVector3(start_x, start_y+1.5f-11.f,start_z)); + } + rb->setCenterOfMassTransform(startTransform); + } + return; + } +#endif +// we don't use motionState, so reset transforms here + int numObjects = m_dynamicsWorld->getNumCollisionObjects(); + btCollisionObjectArray& collisionObjects = m_dynamicsWorld->getCollisionObjectArray(); + + float start_x = START_POS_X - ARRAY_SIZE_X * SCALING; + float start_y = START_POS_Y - ARRAY_SIZE_Y * SCALING; + float start_z = START_POS_Z - ARRAY_SIZE_Z * SCALING; + btTransform startTransform; + startTransform.setIdentity(); + + for(int n = 0; n < numObjects; n++) + { + btCollisionObject* colObj = collisionObjects[n]; + btRigidBody* rb = btRigidBody::upcast(colObj); + int offs = ARRAY_SIZE_X * ARRAY_SIZE_Z; + int indx = n; + int ky = indx / offs; + indx -= ky * offs; + int kx = indx / ARRAY_SIZE_Z; + indx -= kx * ARRAY_SIZE_Z; + int kz = indx; + startTransform.setOrigin(SCALING*btVector3( + 2.0*SCALING*kx + start_x, + 2.0*SCALING*ky + start_y, + 2.0*SCALING*kz + start_z)); + rb->setCenterOfMassTransform(startTransform); + } +} + + + +void BasicDemo3D::exitPhysics() +{ + + //cleanup in the reverse order of creation/initialization + + //remove the rigidbodies from the dynamics world and delete them + int i; + for (i=m_dynamicsWorld->getNumCollisionObjects()-1; i>=0 ;i--) + { + btCollisionObject* obj = m_dynamicsWorld->getCollisionObjectArray()[i]; + btRigidBody* body = btRigidBody::upcast(obj); + if (body && body->getMotionState()) + { + delete body->getMotionState(); + } + m_dynamicsWorld->removeCollisionObject( obj ); + delete obj; + } + + //delete collision shapes + for (int j=0;jgetSimulationIslandManager()->setSplitIslands(sCurrSolverIndex == 0); + pDdw->setConstraintSolver(sConstraintSolvers[sCurrSolverIndex]); + printf("\nUsing %s\n", sConstraintSolverNames[sCurrSolverIndex]); + } + break; +#endif + case 'c' : + { + gbDrawBatches = !gbDrawBatches; + break; + } + case 'b' : + { + gSelectedBatch++; + gSelectedBatch %= (CUDA_DEMO_DYNAMICS_WORLD3D_MAX_BATCHES + 1); + break; + } + case 'u' : + { +#ifdef BT_USE_CUDA + btCudaDemoDynamicsWorld3D* pDdw = (btCudaDemoDynamicsWorld3D*)m_dynamicsWorld; + gUseCPUSolver = !gUseCPUSolver; + pDdw->setUseCPUSolver(gUseCPUSolver); +#endif + break; + } + case 'w' : + { + gDrawWire = !gDrawWire; + setWireMode(gDrawWire); + break; + } + case 'm' : + { + btCudaDemoDynamicsWorld3D* pDdw = (btCudaDemoDynamicsWorld3D*)m_dynamicsWorld; + gUseCudaMotIntegr = !gUseCudaMotIntegr; + pDdw->setUseCudaMotIntegr(gUseCudaMotIntegr); + break; + } + + default : + { + DemoApplication::keyboardCallback(key, x, y); + } + break; + } + + if(key == ' ') + { + //gPairCache->reset(); + } +} + + +void BasicDemo3D::mouseFunc(int button, int state, int x, int y) +{ + if (state == GLUT_DOWN) { + m_mouseButtons |= 1<m_numInBatches[i]); + GLDebugDrawString(xOffs-80, yOffs,buf); + yOffs += 15.f; + } +} + + +void BasicDemo3D::renderme() +{ + updateCamera(); + + glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); + + if(gDrawWire) + { + glColor3f(1.f, 1.f, 1.f); + glDisable(GL_LIGHTING); + setTexturing(false); + } + else + { + myinit(); + setTexturing(true); + } + + renderscene(0); + + if(gbDrawBatches) + { + ((btCudaDemoDynamicsWorld3D*)m_dynamicsWorld)->debugDrawConstraints(gSelectedBatch, cBatchColorTab); + } + glColor3f(0, 0, 0); + if ((m_debugMode & btIDebugDraw::DBG_NoHelpText)==0) + { + setOrthographicProjection(); + int xOffset = 10.f; + int yStart = 20.f; + int yIncr = 20.f; + showProfileInfo(xOffset, yStart, yIncr); + DrawConstraintInfo(); + outputDebugInfo(xOffset, yStart, yIncr); + resetPerspectiveProjection(); + } +} + + + +extern int gNumClampedCcdMotions; +#define SHOW_NUM_DEEP_PENETRATIONS 1 +#ifdef SHOW_NUM_DEEP_PENETRATIONS + extern int gNumDeepPenetrationChecks; + extern int gNumSplitImpulseRecoveries; + extern int gNumGjkChecks; + extern int gNumAlignedAllocs; + extern int gNumAlignedFree; + extern int gTotalBytesAlignedAllocs; +#endif // + + +void BasicDemo3D::outputDebugInfo(int & xOffset,int & yStart, int yIncr) +{ + char buf[124]; + glDisable(GL_LIGHTING); + glColor3f(0, 0, 0); + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"mouse to interact"); + GLDebugDrawString(xOffset,yStart,buf); + yStart += yIncr; + + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"space to reset"); + GLDebugDrawString(xOffset,yStart,buf); + yStart += yIncr; + + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"cursor keys and z,x to navigate"); + GLDebugDrawString(xOffset,yStart,buf); + yStart += yIncr; + + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"i to toggle simulation, s single step"); + GLDebugDrawString(xOffset,yStart,buf); + yStart += yIncr; + + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"q to quit"); + GLDebugDrawString(xOffset,yStart,buf); + yStart += yIncr; + + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"h to toggle help text"); + GLDebugDrawString(xOffset,yStart,buf); + yStart += yIncr; + + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"p to toggle profiling (+results to file)"); + GLDebugDrawString(xOffset,yStart,buf); + yStart += yIncr; + + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"w to toggle wireframe/solid rendering"); + GLDebugDrawString(xOffset,yStart,buf); + yStart += yIncr; + + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"c to toggle constraint drawing"); + GLDebugDrawString(xOffset,yStart,buf); + yStart += yIncr; + + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"b to draw single constraint batch"); + GLDebugDrawString(xOffset,yStart,buf); + yStart += yIncr; + + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"u to toggle between CPU and CUDA solvers"); + GLDebugDrawString(xOffset,yStart,buf); + yStart += yIncr; + + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"d to toggle between different batch builders"); + GLDebugDrawString(xOffset,yStart,buf); + yStart += yIncr; + + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"m to toggle between CUDA / CPU motion integrators"); + GLDebugDrawString(xOffset,yStart,buf); + yStart += yIncr; + + if (getDynamicsWorld()) + { + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"# objects = %d",getDynamicsWorld()->getNumCollisionObjects()); + GLDebugDrawString(xOffset,yStart,buf); + yStart += yIncr; + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"# pairs = %d",getDynamicsWorld()->getBroadphase()->getOverlappingPairCache()->getNumOverlappingPairs()); + GLDebugDrawString(xOffset,yStart,buf); + +/*@@ + + yStart += yIncr; + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"# skipped collisions=%d",gSkippedCol); + GLDebugDrawString(xOffset,yStart,buf); + + yStart += yIncr; + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"# processed collisions=%d",gProcessedCol); + GLDebugDrawString(xOffset,yStart,buf); + + yStart += yIncr; + glRasterPos3f(xOffset,yStart,0); + sprintf(buf,"culled narrowphase collisions=%f",btScalar(gSkippedCol)/(gProcessedCol+gSkippedCol)); + GLDebugDrawString(xOffset,yStart,buf); + yStart += yIncr; +@@*/ + + + + } +} // BasicDemo3D::outputDebugInfo() + +void BasicDemo3D::setWireMode(bool wireOnOff) +{ + int dbgDrawMode = m_dynamicsWorld->getDebugDrawer()->getDebugMode(); + if(wireOnOff) + { + dbgDrawMode |= btIDebugDraw::DBG_FastWireframe; + } + else + { + dbgDrawMode &= ~btIDebugDraw::DBG_FastWireframe; + } + m_dynamicsWorld->getDebugDrawer()->setDebugMode(dbgDrawMode); + m_debugMode = dbgDrawMode; +} // BasicDemo3D::setWireMode() diff --git a/Demos/Gpu3dDemo/BasicDemo3d.h b/Demos/Gpu3dDemo/BasicDemo3d.h new file mode 100644 index 000000000..7fe6625cd --- /dev/null +++ b/Demos/Gpu3dDemo/BasicDemo3d.h @@ -0,0 +1,93 @@ +/* +Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org +Copyright (C) 2006, 2007 Sony Computer Entertainment Inc. + +This software is provided 'as-is', without any express or implied warranty. +In no event will the authors be held liable for any damages arising from the use of this software. +Permission is granted to anyone to use this software for any purpose, +including commercial applications, and to alter it and redistribute it freely, +subject to the following restrictions: + +1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. +2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. +3. This notice may not be removed or altered from any source distribution. +*/ + +#ifndef BASIC_DEMO3D_H +#define BASIC_DEMO3D_H + +#include "GlutDemoApplication.h" +#include "LinearMath/btAlignedObjectArray.h" +#include "BulletDynamics/Dynamics/btDiscreteDynamicsWorld.h" + +class btBroadphaseInterface; +class btCollisionShape; +class btOverlappingPairCache; +class btCollisionDispatcher; +class btConstraintSolver; +struct btCollisionAlgorithmCreateFunc; +class btDefaultCollisionConfiguration; + +///BasicDemo is good starting point for learning the code base and porting. +class BasicDemo3D : public GlutDemoApplication +{ + + //keep the collision shapes, for deletion/cleanup + btAlignedObjectArray m_collisionShapes; + + btBroadphaseInterface* m_broadphase; + + btCollisionDispatcher* m_dispatcher; + + btConstraintSolver* m_solver; + + btDefaultCollisionConfiguration* m_collisionConfiguration; + + int m_mouseButtons; + int m_mouseOldX; + int m_mouseOldY; + + public: + + BasicDemo3D() + { + } + virtual ~BasicDemo3D() + { + exitPhysics(); + } + void initPhysics(); + + void exitPhysics(); + + virtual void clientMoveAndDisplay(); + + virtual void displayCallback(); + + virtual void keyboardCallback(unsigned char key, int x, int y); + virtual void mouseFunc(int button, int state, int x, int y); + virtual void mouseMotionFunc(int x,int y); + + virtual void clientResetScene(); + + static DemoApplication* Create() + { + BasicDemo3D* demo = new BasicDemo3D; + demo->myinit(); + demo->initPhysics(); + demo->m_mouseButtons = 0; + demo->m_mouseOldX = 0; + demo->m_mouseOldY = 0; + return demo; + } + + void DrawConstraintInfo(); + void outputDebugInfo(int & xOffset,int & yStart, int yIncr); + virtual void renderme(); + + void setWireMode(bool wireOnOff); +}; + + +#endif //BASIC_DEMO3D_H + diff --git a/Demos/Gpu3dDemo/CMakeLists.txt b/Demos/Gpu3dDemo/CMakeLists.txt new file mode 100644 index 000000000..2b3093d43 --- /dev/null +++ b/Demos/Gpu3dDemo/CMakeLists.txt @@ -0,0 +1,34 @@ +# This is basically the overall name of the project in Visual Studio this is the name of the Solution File + + +# For every executable you have with a main method you should have an add_executable line below. +# For every add executable line you should list every .cpp and .h file you have associated with that executable. + + +# This is the variable for Windows. I use this to define the root of my directory structure. +SET(GLUT_ROOT ${BULLET_PHYSICS_SOURCE_DIR}/Glut) + +# You shouldn't have to modify anything below this line +######################################################## + +INCLUDE_DIRECTORIES( +${BULLET_PHYSICS_SOURCE_DIR}/src ${BULLET_PHYSICS_SOURCE_DIR}/Demos/OpenGL } +) + +LINK_LIBRARIES( +OpenGLSupport BulletDynamics BulletCollision LinearMath ${GLUT_glut_LIBRARY} ${OPENGL_gl_LIBRARY} ${OPENGL_glu_LIBRARY} +) + +ADD_EXECUTABLE(AppGpu3dDemo + main.cpp + BasicDemo3d.cpp + BasicDemo3d.h + btGpuDemo3dSharedTypes.h + btGpuDemo3dCpuFunc.cpp + btGpuDemoDynamicsWorld3D.cpp + btGpuDemoDynamicsWorld3D.h + btGpuDemo3dSharedCode.h + btGpuDemo3dSharedDefs.h +) + + diff --git a/Demos/Gpu3dDemo/btGpuDemo3dCpuFunc.cpp b/Demos/Gpu3dDemo/btGpuDemo3dCpuFunc.cpp new file mode 100644 index 000000000..fabb17493 --- /dev/null +++ b/Demos/Gpu3dDemo/btGpuDemo3dCpuFunc.cpp @@ -0,0 +1,32 @@ +/* +Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org +Copyright (C) 2006, 2007 Sony Computer Entertainment Inc. + +This software is provided 'as-is', without any express or implied warranty. +In no event will the authors be held liable for any damages arising from the use of this software. +Permission is granted to anyone to use this software for any purpose, +including commercial applications, and to alter it and redistribute it freely, +subject to the following restrictions: + +1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. +2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. +3. This notice may not be removed or altered from any source distribution. +*/ + +//---------------------------------------------------------------------------------------- + +#include "LinearMath/btQuickprof.h" + +//---------------------------------------------------------------------------------------- + +#include "btGpuDemo3dSharedTypes.h" + +//---------------------------------------------------------------------------------------- + +#include "BulletMultiThreaded/btGpuDefines.h" +#include "BulletMultiThreaded/btGpuUtilsSharedDefs.h" +#include "btGpuDemo3dSharedCode.h" + +//-------------------------------------------------------------------------- +//-------------------------------------------------------------------------- +//-------------------------------------------------------------------------- diff --git a/Demos/Gpu3dDemo/btGpuDemo3dSharedCode.h b/Demos/Gpu3dDemo/btGpuDemo3dSharedCode.h new file mode 100644 index 000000000..d7868d968 --- /dev/null +++ b/Demos/Gpu3dDemo/btGpuDemo3dSharedCode.h @@ -0,0 +1,542 @@ +/* +Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org +Copyright (C) 2006, 2007 Sony Computer Entertainment Inc. + +This software is provided 'as-is', without any express or implied warranty. +In no event will the authors be held liable for any damages arising from the use of this software. +Permission is granted to anyone to use this software for any purpose, +including commercial applications, and to alter it and redistribute it freely, +subject to the following restrictions: + +1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. +2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. +3. This notice may not be removed or altered from any source distribution. +*/ + +//---------------------------------------------------------------------------------------- + +#define USE_FRICTION 1 +#define FRICTION_BOX_GROUND_FACT 0.01f +#define FRICTION_BOX_BOX_FACT 0.01f +//#define FRICTION_BOX_BOX_FACT 0.05f +#define USE_CENTERS 1 + +//---------------------------------------------------------------------------------------- +//---------------------------------------------------------------------------------------- +//---------- C o n s t r a i n t s o l v e r d e m o 3D -------------------------- +//---------------------------------------------------------------------------------------- +//---------------------------------------------------------------------------------------- + +// kernel functions + + +BT_GPU___global__ void clearAccumulationOfLambdaDtD(float* lambdaDtBox, int numConstraints, int numContPoints) +{ + int index = BT_GPU___mul24(BT_GPU_blockIdx.x, BT_GPU_blockDim.x) + BT_GPU_threadIdx.x; + if(index < numConstraints) + { + for(int i=0; i < numContPoints; i++) + lambdaDtBox[numContPoints * index + i] = 0; + } +} // clearAccumulationOfLambdaDtD() + +//---------------------------------------------------------------------------------------- + +BT_GPU___device__ float computeImpulse3D(float3 rVel, + float positionConstraint, + float3 cNormal, + float dt) +{ + const float collisionConstant = 0.1f; + const float baumgarteConstant = 0.1f; + const float penetrationError = 0.02f; + + float lambdaDt=0; + float3 impulse=BT_GPU_make_float3(0.f,0.f,0.f); + + if(positionConstraint >= 0) + return lambdaDt; + + positionConstraint = min(0.0f,positionConstraint+penetrationError); + + lambdaDt = -(BT_GPU_dot(cNormal,rVel)*(collisionConstant)); + lambdaDt -= (baumgarteConstant/dt*positionConstraint); + + return lambdaDt; +} // computeImpulse3D() + +//---------------------------------------------------------------------------------------- + +#if 0 +#define VLIM 1000.f +void BT_GPU___device__ chk_vect(float4* v) +{ + if(v->x < -VLIM) v->x = 0.f; + if(v->x > VLIM) v->x = 0.f; + if(v->y < -VLIM) v->y = 0.f; + if(v->y > VLIM) v->y = 0.f; + if(v->z < -VLIM) v->z = 0.f; + if(v->z > VLIM) v->z = 0.f; +} // chk_vect() +#endif + +//---------------------------------------------------------------------------------------- + +BT_GPU___global__ void collisionWithWallBox3DD(float4 *trans, + float4 *vel, + float4* angVel, + btCudaPartProps pProp, + btCudaBoxProps gProp, + int nParticles, + float dt) +{ + int idx = BT_GPU___mul24(BT_GPU_blockIdx.x, BT_GPU_blockDim.x) + BT_GPU_threadIdx.x; + float3 aPos; + float positionConstraint; + float3 impulse; + + if(idx < nParticles) + { + aPos = BT_GPU_make_float34(trans[idx * 4 + 3]); + for(int iVtx=0; iVtx < 8; iVtx++) + { + float3 dx = BT_GPU_make_float34(trans[idx * 4 + 0]); + float3 dy = BT_GPU_make_float34(trans[idx * 4 + 1]); + float3 dz = BT_GPU_make_float34(trans[idx * 4 + 2]); + float3 rerVertex = ((iVtx & 1) ? dx : dx * (-1.f)); + rerVertex += ((iVtx & 2) ? dy : dy * (-1.f)); + rerVertex += ((iVtx & 4) ? dz : dz * (-1.f)); + float3 vPos = aPos + rerVertex; + float3 aVel = BT_GPU_make_float3(vel[idx].x, vel[idx].y, vel[idx].z); + float3 aAngVel = BT_GPU_make_float34(angVel[idx]); + float3 vVel =aVel+BT_GPU_cross(aAngVel, rerVertex); + float restitution=0.5; + { + positionConstraint = vPos.y - gProp.minY; + impulse = BT_GPU_make_float31(0.0f); + if(positionConstraint < 0) + { + float3 groundNormal; + groundNormal = BT_GPU_make_float3(0.0f,1.0f,0.0f); + impulse = groundNormal * restitution * computeImpulse3D(vVel, positionConstraint, groundNormal, dt); +#if USE_FRICTION // only with ground for now + float3 lat_vel = vVel - groundNormal * BT_GPU_dot(groundNormal,vVel); + float lat_vel_len = BT_GPU_dot(lat_vel, lat_vel); + if (lat_vel_len > 0) + { + lat_vel_len = sqrtf(lat_vel_len); + lat_vel *= 1.f/lat_vel_len; + impulse -= lat_vel * BT_GPU_dot(lat_vel, vVel) * FRICTION_BOX_GROUND_FACT; + } +#endif //USE_FRICTION + vel[idx] += BT_GPU_make_float42(impulse,0.0f); + angVel[idx] += BT_GPU_make_float42(BT_GPU_cross(rerVertex,impulse), 0.0f); + } + } + { + positionConstraint = vPos.x - gProp.minX; + impulse = BT_GPU_make_float31(0.0f); + if(positionConstraint < 0) + { + float3 normal = BT_GPU_make_float3(1.0f,0.0f,0.0f); + impulse = normal * restitution * computeImpulse3D(vVel,positionConstraint,normal,dt); + vel[idx] += BT_GPU_make_float42(impulse,0.0f); + angVel[idx] += BT_GPU_make_float42(BT_GPU_cross(rerVertex,impulse), 0.0f); + } + } + { + positionConstraint = gProp.maxX - vPos.x; + impulse = BT_GPU_make_float31(0.0f); + if(positionConstraint < 0) + { + float3 normal = BT_GPU_make_float3(-1.0f,0.0f,0.0f); + impulse = normal * restitution * computeImpulse3D(vVel,positionConstraint,normal,dt); + vel[idx] += BT_GPU_make_float42(impulse,0.0f); + angVel[idx] += BT_GPU_make_float42(BT_GPU_cross(rerVertex,impulse), 0.0f); + } + } + { + positionConstraint = vPos.z - gProp.minZ; + impulse = BT_GPU_make_float31(0.0f); + if(positionConstraint < 0) + { + float3 normal = BT_GPU_make_float3(0.0f,0.0f,1.0f); + impulse = normal * restitution * computeImpulse3D(vVel,positionConstraint,normal,dt); + vel[idx] += BT_GPU_make_float42(impulse,0.0f); + angVel[idx] += BT_GPU_make_float42(BT_GPU_cross(rerVertex,impulse), 0.0f); + } + } + { + positionConstraint = gProp.maxZ - vPos.z; + impulse = BT_GPU_make_float31(0.0f); + if(positionConstraint < 0) + { + float3 normal = BT_GPU_make_float3(0.0f,0.0f,-1.0f); + impulse = normal * restitution * computeImpulse3D(vVel,positionConstraint,normal,dt); + vel[idx] += BT_GPU_make_float42(impulse,0.0f); + angVel[idx] += BT_GPU_make_float42(BT_GPU_cross(rerVertex,impulse), 0.0f); + } + } + } + } +} // collisionWithWallBox3DD() + +//---------------------------------------------------------------------------------------- + +BT_GPU___global__ void collisionBatchResolutionBox3DD(int2 *constraints, + int *batch, + int nConstraints, + float4 *trans, + float4 *vel, + float4 *angularVel, + float *lambdaDtBox, + float *iPositionConstraint, + float3 *normal, + float3 *contact, + btCudaPartProps pProp, + int iBatch, + float dt) +{ + float3 relVel; + float3 impulse; + float lambdaDt; + float positionConstraint; + int k_idx = BT_GPU___mul24(BT_GPU_blockIdx.x, BT_GPU_blockDim.x) + BT_GPU_threadIdx.x; + if(k_idx < nConstraints) + { + int idx = batch[k_idx]; + int aId=constraints[idx].x; + int bId=constraints[idx].y; + float3 aPos = BT_GPU_make_float34(trans[aId * 4 + 3]); + float3 bPos = BT_GPU_make_float34(trans[bId * 4 + 3]); + float3 aVel = BT_GPU_make_float34(vel[aId]); + float3 bVel = BT_GPU_make_float34(vel[bId]); + float3 aAngVel = BT_GPU_make_float34(angularVel[aId]); + float3 bAngVel = BT_GPU_make_float34(angularVel[bId]); + for(int iVtx = 0; iVtx < 4; iVtx++) + { + float3 contactPoint = contact[idx * 4 + iVtx] - aPos; + positionConstraint = iPositionConstraint[idx * 4 + iVtx]; + if(positionConstraint > 0) + { + float3 contactNormal = normal[idx * 4 + iVtx]; + relVel = (aVel + BT_GPU_cross(aAngVel, contactPoint)) + -(bVel + BT_GPU_cross(bAngVel, contactPoint+aPos-bPos)); + + lambdaDt= computeImpulse3D(relVel, -positionConstraint, contactNormal, dt); + { + float rLambdaDt=lambdaDtBox[idx * 4 + iVtx]; + float pLambdaDt=rLambdaDt; + rLambdaDt=max(pLambdaDt+lambdaDt,0.0f); + lambdaDt=rLambdaDt-pLambdaDt; + lambdaDtBox[idx * 4 + iVtx]=rLambdaDt; + } + impulse = contactNormal*lambdaDt*0.5; +#if USE_FRICTION + float3 lat_vel = relVel - contactNormal * BT_GPU_dot(contactNormal, relVel); + float lat_vel_len = BT_GPU_dot(lat_vel, lat_vel); + if (lat_vel_len > 0) + { + lat_vel_len = sqrtf(lat_vel_len); + lat_vel *= 1.f/lat_vel_len; + impulse -= lat_vel * BT_GPU_dot(lat_vel , relVel) * FRICTION_BOX_BOX_FACT; + } +#endif //USE_FRICTION + aVel+= impulse; + bVel-= impulse; + aAngVel += BT_GPU_cross(contactPoint, impulse); + bAngVel -= BT_GPU_cross(contactPoint+aPos-bPos, impulse); + } + } + vel[aId]=BT_GPU_make_float42(aVel,0.0f); + vel[bId]=BT_GPU_make_float42(bVel,0.0f); + angularVel[aId]=BT_GPU_make_float42(aAngVel,0.0f); + angularVel[bId]=BT_GPU_make_float42(bAngVel,0.0f); + } +} // collisionBatchResolutionBox3DD() + +//---------------------------------------------------------------------------------------- +//---------------------------------------------------------------------------------------- +//---------------------------------------------------------------------------------------- + + +extern "C" +{ + +// global functions + +//---------------------------------------------------------------------------------------- + +void BT_GPU_PREF(clearAccumulationOfLambdaDt(float* lambdaDtBox, int numConstraints, int numContPoints)) +{ + if(!numConstraints) + { + return; + } + int numThreads, numBlocks; + BT_GPU_PREF(computeGridSize)(numConstraints, 256, numBlocks, numThreads); + // execute the kernel + BT_GPU_EXECKERNEL(numBlocks, numThreads, clearAccumulationOfLambdaDtD, (lambdaDtBox, numConstraints, numContPoints)); + // check if kernel invocation generated an error + BT_GPU_CHECK_ERROR("clearAccumulationOfLambdaDtD kernel execution failed"); + +} // clearAccumulationOfLambdaDt() + +//---------------------------------------------------------------------------------------- + +void BT_GPU_PREF(collisionWithWallBox3D(void* trans,void* vel,void* angVel,btCudaPartProps pProp, btCudaBoxProps gProp,int numObjs,float dt)) +{ + if(!numObjs) + { + return; + } + float4* pTrans = (float4*)trans; + float4* pVel = (float4*)vel; + float4* pAngVel = (float4*)angVel; + int numThreads, numBlocks; + BT_GPU_PREF(computeGridSize)(numObjs, 256, numBlocks, numThreads); + // execute the kernel + BT_GPU_EXECKERNEL(numBlocks, numThreads, collisionWithWallBox3DD, (pTrans,pVel,pAngVel,pProp,gProp,numObjs,dt)); + // check if kernel invocation generated an error + BT_GPU_CHECK_ERROR("collisionWithWallBox3DD kernel execution failed"); +} // collisionWithWallBox3D() + +//---------------------------------------------------------------------------------------- + +void BT_GPU_PREF(collisionBatchResolutionBox3D(void* constraints,int *batch,int numConstraints,void *trans,void *vel, + void *angularVel,float *lambdaDtBox,float *positionConstraint,void* normal,void* contact, + btCudaPartProps pProp,int iBatch,float dt)) +{ + if(!numConstraints) + { + return; + } + int2* pConstr = (int2*)constraints; + float4* pTrans = (float4*)trans; + float4* pVel = (float4*)vel; + float4* pAngVel = (float4*)angularVel; + float3* pNorm = (float3*)normal; + float3* pContact = (float3*)contact; + int numThreads, numBlocks; + BT_GPU_PREF(computeGridSize)(numConstraints, 128, numBlocks, numThreads); + // execute the kernel + BT_GPU_EXECKERNEL(numBlocks, numThreads, collisionBatchResolutionBox3DD, (pConstr,batch,numConstraints,pTrans,pVel,pAngVel,lambdaDtBox,positionConstraint,pNorm,pContact,pProp,iBatch,dt)); + // check if kernel invocation generated an error + BT_GPU_CHECK_ERROR("collisionBatchResolutionBox3DD kernel execution failed"); +} // collisionBatchResolutionBox3D() + +//---------------------------------------------------------------------------------------- + +} // extern "C" + +//---------------------------------------------------------------------------------------- +//---------------------------------------------------------------------------------------- +//---------- M o t i o n i n t e g r a t o r d e m o ----------------------------- +//---------------------------------------------------------------------------------------- +//---------------------------------------------------------------------------------------- + +// kernel functions + +BT_GPU___global__ void integrVelD(float4* pForceTorqueDamp, float4* pInvInertiaMass, float4* pVel, float4* pAngVel, float timeStep, unsigned int numBodies) +{ + int index = BT_GPU___mul24(BT_GPU_blockIdx.x, BT_GPU_blockDim.x) + BT_GPU_threadIdx.x; + if(index >= (int)numBodies) + { + return; + } + // unpack input data + float3 force = BT_GPU_make_float34(pForceTorqueDamp[index * 2]); + float lin_damp = pForceTorqueDamp[index * 2].w; + float3 torque = BT_GPU_make_float34(pForceTorqueDamp[index * 2 + 1]); + float ang_damp = pForceTorqueDamp[index * 2 + 1].w; + float3 linVel = BT_GPU_make_float34(pVel[index]); + float3 angVel = BT_GPU_make_float34(pAngVel[index]); + float3 in_mass_0 = BT_GPU_make_float34(pInvInertiaMass[index * 3]); + float3 in_mass_1 = BT_GPU_make_float34(pInvInertiaMass[index * 3 + 1]); + float3 in_mass_2 = BT_GPU_make_float34(pInvInertiaMass[index * 3 + 2]); + float mass = pInvInertiaMass[index * 3].w; + // integrate linear velocity + float3 outLinVel, outAngVel; + outLinVel = linVel + force * mass * timeStep; + // integrate angular velocity + outAngVel.x = BT_GPU_dot(in_mass_0, torque); + outAngVel.y = BT_GPU_dot(in_mass_1, torque); + outAngVel.z = BT_GPU_dot(in_mass_2, torque); + outAngVel += angVel; + /// clamp angular velocity. collision calculations will fail on higher angular velocities + #if(!defined(M_PI)) + #define M_PI 3.1415926f + #endif + #define BT_CUDA_MAX_SQ_ANGVEL (M_PI*M_PI) + float sq_angvel = BT_GPU_dot(outAngVel, outAngVel); + sq_angvel *= timeStep * timeStep; + float fact; + if(sq_angvel > BT_CUDA_MAX_SQ_ANGVEL) + { + fact = sqrtf(BT_CUDA_MAX_SQ_ANGVEL/sq_angvel) / timeStep; + outAngVel *= fact; + } + // now apply damping + fact = powf(1.0f - lin_damp, timeStep); + outLinVel *= fact; + fact = powf(1.0f - ang_damp, timeStep); + outAngVel *= fact; + // pack results + pVel[index] = BT_GPU_make_float42(outLinVel, 0.f); + pAngVel[index] = BT_GPU_make_float42(outAngVel, 0.f); +} // integrVelD() + +#define BT_GPU__ANGULAR_MOTION_THRESHOLD (0.25f * M_PI) + +//---------------------------------------------------------------------------------------- + +BT_GPU___device__ float4 getRotation(float4* trans) +{ + float trace = trans[0].x + trans[1].y + trans[2].z; + float temp[4]; + if(trace > 0.0f) + { + float s = sqrtf(trace + 1.0f); + temp[3] = s * 0.5f; + s = 0.5f / s; + temp[0] = (trans[1].z - trans[2].y) * s; + temp[1] = (trans[2].x - trans[0].z) * s; + temp[2] = (trans[0].y - trans[1].x) * s; + } + else + { + typedef float btMatrRow[4]; + btMatrRow* m_el = (btMatrRow*)trans; + int i = m_el[0][0] < m_el[1][1] ? + (m_el[1][1] < m_el[2][2] ? 2 : 1) : + (m_el[0][0] < m_el[2][2] ? 2 : 0); + int j = (i + 1) % 3; + int k = (i + 2) % 3; + float s = sqrtf(m_el[i][i] - m_el[j][j] - m_el[k][k] + 1.0f); + temp[i] = s * 0.5f; + s = 0.5f / s; + temp[3] = (m_el[j][k] - m_el[k][j]) * s; + temp[j] = (m_el[i][j] + m_el[j][i]) * s; + temp[k] = (m_el[i][k] + m_el[k][i]) * s; + } + float4 q = BT_GPU_make_float44(temp[0],temp[1],temp[2],temp[3]); + return q; +} // getRotation() + +//---------------------------------------------------------------------------------------- + +BT_GPU___device__ float4 quatMult(float4& q1, float4& q2) +{ + return BT_GPU_make_float44(q1.w * q2.x + q1.x * q2.w + q1.y * q2.z - q1.z * q2.y, + q1.w * q2.y + q1.y * q2.w + q1.z * q2.x - q1.x * q2.z, + q1.w * q2.z + q1.z * q2.w + q1.x * q2.y - q1.y * q2.x, + q1.w * q2.w - q1.x * q2.x - q1.y * q2.y - q1.z * q2.z); +} // quatMult() + +//---------------------------------------------------------------------------------------- + +BT_GPU___device__ void quatNorm(float4& q) +{ + float len = sqrtf(BT_GPU_dot4(q, q)); + q *= 1.f / len; +} // quatNorm() + +//---------------------------------------------------------------------------------------- + +BT_GPU___device__ void setRotation(float4& q, float4* trans) +{ + float d = BT_GPU_dot4(q, q); + float s = 2.0f / d; + float xs = q.x * s, ys = q.y * s, zs = q.z * s; + float wx = q.w * xs, wy = q.w * ys, wz = q.w * zs; + float xx = q.x * xs, xy = q.x * ys, xz = q.x * zs; + float yy = q.y * ys, yz = q.y * zs, zz = q.z * zs; + trans[0].x = 1.0f - (yy + zz); + trans[1].x = xy - wz; + trans[2].x = xz + wy; + trans[0].y = xy + wz; + trans[1].y = 1.0f - (xx + zz); + trans[2].y = yz - wx; + trans[0].z = xz - wy; + trans[1].z = yz + wx; + trans[2].z = 1.0f - (xx + yy); + trans[0].w = trans[1].w = trans[2].w = 0.0f; +} // setRotation() + +//---------------------------------------------------------------------------------------- + +BT_GPU___global__ void integrTransD(float4* pTrans, float4* pVel, float4* pAngVel, float timeStep, unsigned int numBodies) +{ + int index = BT_GPU___mul24(BT_GPU_blockIdx.x, BT_GPU_blockDim.x) + BT_GPU_threadIdx.x; + if(index >= (int)numBodies) + { + return; + } + float3 pos = BT_GPU_make_float34(pTrans[index * 4 + 3]); + float3 linvel = BT_GPU_make_float34(pVel[index]); + pos += linvel * timeStep; + + float3 axis; + float3 angvel = BT_GPU_make_float34(pAngVel[index]); + float fAngle = sqrtf(BT_GPU_dot(angvel, angvel)); + //limit the angular motion + if(fAngle*timeStep > BT_GPU__ANGULAR_MOTION_THRESHOLD) + { + fAngle = BT_GPU__ANGULAR_MOTION_THRESHOLD / timeStep; + } + if(fAngle < 0.001f) + { + // use Taylor's expansions of sync function + axis = angvel * (0.5f*timeStep-(timeStep*timeStep*timeStep)*0.020833333333f * fAngle * fAngle); + } + else + { + // sync(fAngle) = sin(c*fAngle)/t + axis = angvel * ( sinf(0.5f * fAngle * timeStep) / fAngle); + } + float4 dorn = BT_GPU_make_float42(axis, cosf(fAngle * timeStep * 0.5f)); + float4 orn0 = getRotation(pTrans + index * 4); + float4 predictedOrn = quatMult(dorn, orn0); + quatNorm(predictedOrn); + setRotation(predictedOrn, pTrans + index * 4); + pTrans[index * 4 + 3] = BT_GPU_make_float42(pos, 0.f); +} // integrTransD() + + +//---------------------------------------------------------------------------------------- +//---------------------------------------------------------------------------------------- +//---------------------------------------------------------------------------------------- + +// global functions + +extern "C" +{ + +//---------------------------------------------------------------------------------------- + +void BT_GPU_PREF(integrVel(float* pForceTorqueDamp, float* pInvInertiaMass, void* pVel, void* pAngVel, float timeStep, unsigned int numBodies)) +{ + int numThreads, numBlocks; + BT_GPU_PREF(computeGridSize)(numBodies, 256, numBlocks, numThreads); + BT_GPU_EXECKERNEL(numBlocks, numThreads, integrVelD, ((float4*)pForceTorqueDamp, (float4*)pInvInertiaMass, (float4*)pVel, (float4*)pAngVel, timeStep, numBodies)); + BT_GPU_CHECK_ERROR("Kernel execution failed: btCuda_integrVelD"); +} // integrVel() + +//---------------------------------------------------------------------------------------- + +void BT_GPU_PREF(integrTrans(void* trans, void* vel, void* angVel, float timeStep, int numBodies)) +{ + int numThreads, numBlocks; + BT_GPU_PREF(computeGridSize)(numBodies, 256, numBlocks, numThreads); + BT_GPU_EXECKERNEL(numBlocks, numThreads, integrTransD, ((float4*)trans, (float4*)vel, (float4*)angVel, timeStep, numBodies)); + BT_GPU_CHECK_ERROR("Kernel execution failed: btCuda_integrTransD"); +} // integrTrans() + +//---------------------------------------------------------------------------------------- + +} // extern "C" + +//------------------------------------------------------------------------------------------------ +//------------------------------------------------------------------------------------------------ +//------------------------------------------------------------------------------------------------ +//---------------------------------------------------------------------------------------- diff --git a/Demos/Gpu3dDemo/btGpuDemo3dSharedDefs.h b/Demos/Gpu3dDemo/btGpuDemo3dSharedDefs.h new file mode 100644 index 000000000..6e4adaa8a --- /dev/null +++ b/Demos/Gpu3dDemo/btGpuDemo3dSharedDefs.h @@ -0,0 +1,38 @@ +/* +Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org +Copyright (C) 2006, 2007 Sony Computer Entertainment Inc. + +This software is provided 'as-is', without any express or implied warranty. +In no event will the authors be held liable for any damages arising from the use of this software. +Permission is granted to anyone to use this software for any purpose, +including commercial applications, and to alter it and redistribute it freely, +subject to the following restrictions: + +1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. +2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. +3. This notice may not be removed or altered from any source distribution. +*/ + +//---------------------------------------------------------------------------------------- +//---------------------------------------------------------------------------------------- +//---------- C o n s t r a i n t s o l v e r d e m o ---------------------------- +//---------------------------------------------------------------------------------------- +//---------------------------------------------------------------------------------------- + + +extern "C" +{ + +void BT_GPU_PREF(clearAccumulationOfLambdaDt(float* lambdaDtBox, int numConstraints, int numContPoints)); +void BT_GPU_PREF(collisionWithWallBox3D(void* trans,void* vel,void* angVel,btCudaPartProps pProp, btCudaBoxProps gProp,int numObjs,float dt)); +void BT_GPU_PREF(collisionBatchResolutionBox3D(void* constraints,int *batch,int numConstraints,void *trans,void *vel, + void *angularVel,float *lambdaDtBox,float *positionConstraint,void* normal,void* contact, + btCudaPartProps pProp,int iBatch,float dt)); + +void BT_GPU_PREF(integrVel(float* pForceTorqueDamp, float* pInvInertiaMass, void* pVel, void* pAngVel, float timeStep, unsigned int numBodies)); +void BT_GPU_PREF(integrTrans(void* trans, void* vel, void* angVel, float timeStep, int numBodies)); + + +} // extern "C" + +//---------------------------------------------------------------------------------------- diff --git a/Demos/Gpu3dDemo/btGpuDemo3dSharedTypes.h b/Demos/Gpu3dDemo/btGpuDemo3dSharedTypes.h new file mode 100644 index 000000000..35713a22c --- /dev/null +++ b/Demos/Gpu3dDemo/btGpuDemo3dSharedTypes.h @@ -0,0 +1,39 @@ +/* +Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org +Copyright (C) 2006, 2007 Sony Computer Entertainment Inc. + +This software is provided 'as-is', without any express or implied warranty. +In no event will the authors be held liable for any damages arising from the use of this software. +Permission is granted to anyone to use this software for any purpose, +including commercial applications, and to alter it and redistribute it freely, +subject to the following restrictions: + +1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. +2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. +3. This notice may not be removed or altered from any source distribution. +*/ + +//---------------------------------------------------------------------------------------- +//---------------------------------------------------------------------------------------- +//---------- C o n s t r a i n t s o l v e r d e m o ---------------------------- +//---------------------------------------------------------------------------------------- +//---------------------------------------------------------------------------------------- + +struct btCudaPartProps +{ + float m_mass; + float m_diameter; + float m_restCoeff; +}; + +struct btCudaBoxProps +{ + float minX; + float maxX; + float minY; + float maxY; + float minZ; + float maxZ; +}; + +//---------------------------------------------------------------------------------------- diff --git a/Demos/Gpu3dDemo/btGpuDemoDynamicsWorld3D.cpp b/Demos/Gpu3dDemo/btGpuDemoDynamicsWorld3D.cpp new file mode 100644 index 000000000..ad014cd86 --- /dev/null +++ b/Demos/Gpu3dDemo/btGpuDemoDynamicsWorld3D.cpp @@ -0,0 +1,593 @@ +/* +Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org +Copyright (C) 2006, 2007 Sony Computer Entertainment Inc. + +This software is provided 'as-is', without any express or implied warranty. +In no event will the authors be held liable for any damages arising from the use of this software. +Permission is granted to anyone to use this software for any purpose, +including commercial applications, and to alter it and redistribute it freely, +subject to the following restrictions: + +1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. +2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. +3. This notice may not be removed or altered from any source distribution. +*/ + +//-------------------------------------------------------------------------- + +#include "btGpuDemoDynamicsWorld3D.h" +#include "BulletCollision/CollisionDispatch/btCollisionDispatcher.h" +#include "BulletCollision/BroadphaseCollision/btSimpleBroadphase.h" +#include "BulletCollision/CollisionShapes/btCollisionShape.h" +#include "BulletDynamics/Dynamics/btRigidBody.h" +#include "BulletDynamics/ConstraintSolver/btSequentialImpulseConstraintSolver.h" +#include "BulletDynamics/ConstraintSolver/btContactSolverInfo.h" +#include "LinearMath/btQuickprof.h" +#include "GlutStuff.h" + +#include + +//-------------------------------------------------------------------------- + +#define BT_GPU_PREF(func) btCuda_##func +#include "../../src/BulletMultiThreaded/btGpuUtilsSharedDefs.h" +#include "btGpuDemo3dSharedDefs.h" +#undef BT_GPU_PREF + +#define BT_GPU_PREF(func) btGpu_##func +#include "btGpuDemo3dSharedDefs.h" +#undef BT_GPU_PREF + +//-------------------------------------------------------------------------- + +#if 0 +static void check_vel(btVector3& v, int id, char* tag) +{ + int i; + for(i = 0; i < 3; i++) + { + btScalar a = v[i]; + a = btFabs(a); + if(a > 1000.f) + { + break; + } + } + if(i < 3) + { + printf("\nERROR in %s (%4d) : %7.2f %7.2f %7.2f\n", tag, id, v[0], v[1], v[2]); + v[0] = v[1] = v[2] = 0.f; + } +} +#endif + +//-------------------------------------------------------------------------- + +void btCudaDemoDynamicsWorld3D::grabObjData() +{ + int i; + m_numObj = getNumCollisionObjects(); + for(i = 0; i < m_numObj; i++) + { + btCollisionObject* colObj = m_collisionObjects[i]; + colObj->setCompanionId(i); + btRigidBody* rb = btRigidBody::upcast(colObj); + btVector3 v; + if(m_copyIntegrDataToGPU) + { + const btTransform& tr = rb->getCenterOfMassTransform(); + v = tr.getBasis().getColumn(0); + m_hTrans[i * 4 + 0] = *((float4*)&v); + v = tr.getBasis().getColumn(1); + m_hTrans[i * 4 + 1] = *((float4*)&v); + v = tr.getBasis().getColumn(2); + m_hTrans[i * 4 + 2] = *((float4*)&v); + v = rb->getCenterOfMassPosition(); + m_hTrans[i * 4 + 3] = *((float4*)&v); + } + if(!m_useCudaMotIntegr) + { + v = rb->getLinearVelocity(); + m_hVel[i] = *((float4*)&v); + v = rb->getAngularVelocity(); + m_hAngVel[i] = *((float4*)&v); + } + } +} // btCudaDemoDynamicsWorld3D::grabObjData() + +//-------------------------------------------------------------------------- + +void btCudaDemoDynamicsWorld3D::grabConstrData() +{ + int i; + btDispatcher* dispatcher = getDispatcher(); + btPersistentManifold** manifoldPtr = dispatcher->getInternalManifoldPointer(); + int numManifolds = dispatcher->getNumManifolds(); + btPersistentManifold* manifold = 0; + m_numConstraints = 0; + +/* // paranoia + for(int j = 0; j < m_numObj; j++) + { + m_hConstraintCounter[j] = 0; + } +*/ + for(i = 0; i < numManifolds; i++) + { + manifold = manifoldPtr[i]; + int numPoints = manifold->getNumContacts(); + if(!numPoints) + { + continue; + } + + int numActualPoints = 0; + for(int n = 0; n < numPoints; n++) + { + btManifoldPoint& cp = manifold->getContactPoint(n); + if (cp.m_distance1<=0) + { + numActualPoints++; + } + + } + if (!numActualPoints) + continue; + + btRigidBody *rbA, *rbB; + rbA = (btRigidBody*)manifold->getBody0(); + rbB = (btRigidBody*)manifold->getBody1(); + int idA = rbA->getCompanionId(); + int idB = rbB->getCompanionId(); + m_hConstraintCounter[idA]++; + m_hConstraintCounter[idB]++; + if(idA < idB) + { + m_hIds[m_numConstraints].x = idA; + m_hIds[m_numConstraints].y = idB; + + for(int n = 0; n < numPoints; n++) + { + btManifoldPoint& cp = manifold->getContactPoint(n); + btVector3 v = cp.getPositionWorldOnA(); + m_hContact[m_numConstraints * m_maxPointsPerConstr + n] = *((float3*)&v); + v = cp.m_normalWorldOnB; + m_hNormal[m_numConstraints * m_maxPointsPerConstr + n] = *((float3*)&v); + float dist = cp.getDistance(); + if(dist > 0.f) + { + dist = 0.f; + } + m_hPositionConstraint[m_numConstraints * m_maxPointsPerConstr + n] = -dist; + } + } + else + { // should never happen + btAssert(0); + } + for(int n = numPoints; n < m_maxPointsPerConstr; n++) + { + m_hPositionConstraint[m_numConstraints * m_maxPointsPerConstr + n] = 0.f; + } + m_numConstraints++; + } +/* + // paranoia + for(int j = 0; j < m_numObj; j++) + { + if(m_hConstraintCounter[j] > m_maxNeihbors) + { + printf("WARN : constraint connter is %d for object %d\n", m_hConstraintCounter[j], j); + } + } +*/ +} // btCudaDemoDynamicsWorld3D::grabConstrData() + +//-------------------------------------------------------------------------- + +void btCudaDemoDynamicsWorld3D::grabData() +{ + BT_PROFILE("grab data from rigidbody and manifold"); + grabObjData(); + // constraints + grabConstrData(); +} // btCudaDemoDynamicsWorld3D::grabGata() + +//-------------------------------------------------------------------------- + +void btCudaDemoDynamicsWorld3D::createBatches() +{ + BT_PROFILE("create batches"); + int sz = m_numConstraints; + for(int i = 0; i < m_numConstraints; i++) + { + m_hBatchIds[i] = -1; + m_hConstraintUsed[i] = 0; + } + int curBatchId=0; + int* pBatchIds = m_hBatchIds; + int stage; + for(stage = 0; stage < m_maxBatches; stage++) + { // don't print junk on demo screen :-) + m_numInBatches[stage] = 0; + } + for(stage = 0; stage < m_maxBatches; stage++) + { + bool isLast = (stage == m_maxBatches-1); + for(int j = 0; j < m_numObj; j++) + { + m_hConstraintCounter[j] = 0; + } + int numInBatch = 0; + for(int i = 0; i < m_numConstraints; i++) + { + if(m_hConstraintUsed[i]) + { + continue; + } + int2 ids = m_hIds[i]; + if(!isLast) + { + if((m_hConstraintCounter[ids.x] == 0) && (m_hConstraintCounter[ids.y] == 0)) + { + m_hConstraintCounter[ids.x]=1; + m_hConstraintCounter[ids.y]=1; + pBatchIds[numInBatch]=i; + numInBatch++; + m_hConstraintUsed[i] = 1; + } + } + else + { + pBatchIds[numInBatch]=i; + numInBatch++; + m_hConstraintUsed[i] = 1; + } + } + m_numInBatches[stage] = numInBatch; + pBatchIds += numInBatch; + if(!numInBatch) break; + } +} // btCudaDemoDynamicsWorld3D::createBatches() + +//-------------------------------------------------------------------------- + + +void btCudaDemoDynamicsWorld3D::writebackData() +{ + BT_PROFILE("copy velocity into btRigidBody"); + for(int i = 0; i < m_numObj; i++) + { + btCollisionObject* colObj = m_collisionObjects[i]; + btRigidBody* rb = btRigidBody::upcast(colObj); + btVector3 v; + v = *((btVector3*)(m_hVel + i)); + rb->setLinearVelocity(v); + v = *((btVector3*)(m_hAngVel + i)); + rb->setAngularVelocity(v); + } +} // btCudaDemoDynamicsWorld3D::writebackData() + +//-------------------------------------------------------------------------- + +void btCudaDemoDynamicsWorld3D::copyDataToGPU() +{ + BT_PROFILE("copyDataToGPU"); +#ifdef BT_USE_CUDA + btCuda_copyArrayToDevice(m_dIds, m_hIds, sizeof(int2) * m_numConstraints); + btCuda_copyArrayToDevice(m_dBatchIds, m_hBatchIds, sizeof(int) * m_numConstraints); + btCuda_copyArrayToDevice(m_dContact, m_hContact, m_numConstraints * m_maxPointsPerConstr * sizeof(float3)); + btCuda_copyArrayToDevice(m_dNormal, m_hNormal, m_numConstraints * m_maxPointsPerConstr * sizeof(float3)); + btCuda_copyArrayToDevice(m_dPositionConstraint, m_hPositionConstraint, m_numConstraints * m_maxPointsPerConstr * sizeof(float)); + + if(m_copyIntegrDataToGPU) + { + btCuda_copyArrayToDevice(m_dTrans, m_hTrans, m_numObj * sizeof(float4) * 4); + if(m_useCudaMotIntegr) + { + m_copyIntegrDataToGPU = false; + } + } + + if(!m_useCudaMotIntegr) + { + btCuda_copyArrayToDevice(m_dVel, m_hVel, m_numObj * sizeof(float4)); + btCuda_copyArrayToDevice(m_dAngVel, m_hAngVel, m_numObj * sizeof(float4)); + } +#endif +} // btCudaDemoDynamicsWorld3D::copyDataToGPU() + +//-------------------------------------------------------------------------- + +void btCudaDemoDynamicsWorld3D::copyDataFromGPU() +{ + BT_PROFILE("copy velocity data from device"); +#ifdef BT_USE_CUDA + btCuda_copyArrayFromDevice(m_hVel, m_dVel, m_numObj * sizeof(float4)); + btCuda_copyArrayFromDevice(m_hAngVel, m_dAngVel, m_numObj * sizeof(float4)); +#endif +} // btCudaDemoDynamicsWorld3D::copyDataFromGPU() + +//-------------------------------------------------------------------------- + +void btCudaDemoDynamicsWorld3D::solveConstraints(btContactSolverInfo& solverInfo) +{ + if(m_useSeqImpSolver) + { + btDiscreteDynamicsWorld::solveConstraints(solverInfo); + return; + } + if(m_useCPUSolver) + { + solveConstraintsCPU(solverInfo); + return; + } +#ifdef BT_USE_CUDA + BT_PROFILE("solveConstraints"); + grabData(); + createBatches(); + copyDataToGPU(); + + btCudaPartProps partProps; + partProps.m_mass = 1.0f; + partProps.m_diameter = m_objRad * 2.0f; + partProps.m_restCoeff = 1.0f; + + btCudaBoxProps boxProps; + boxProps.minX = m_worldMin[0]; + boxProps.maxX = m_worldMax[0]; + boxProps.minY = m_worldMin[1]; + boxProps.maxY = m_worldMax[1]; + boxProps.minZ = m_worldMin[2]; + boxProps.maxZ = m_worldMax[2]; + { + BT_PROFILE("btCuda_collisionBatchResolutionBox"); + + int nIter=getSolverInfo().m_numIterations; + btDispatcherInfo& dispatchInfo = getDispatchInfo(); + btScalar timeStep = dispatchInfo.m_timeStep; + + btCuda_clearAccumulationOfLambdaDt(m_dLambdaDtBox, m_numConstraints, m_maxPointsPerConstr); + + for(int i=0;igetWorldTransform().getOrigin(); + btVector3 vB = colObjB->getWorldTransform().getOrigin(); + glVertex3f(vA[0], vA[1], vA[2]); + glVertex3f(vB[0], vB[1], vB[2]); + } + pBatchIds += numConstraints; + glEnd(); + } +} // btCudaDemoDynamicsWorld3D::debugDrawConstraints() + +//-------------------------------------------------------------------------- + +void btCudaDemoDynamicsWorld3D::predictUnconstraintMotion(btScalar timeStep) +{ + if(m_useCudaMotIntegr) + { + BT_PROFILE("motIntegr -- predictUnconstraintMotion"); + int i; + { + m_numObj = getNumCollisionObjects(); + float* p_fbuf = m_hForceTorqueDamp; + float* p_mbuf = m_hInvInertiaMass; + for(i = 0; i < m_numObj; i++) + { + btCollisionObject* colObj = m_collisionObjects[i]; + btRigidBody* rb = btRigidBody::upcast(colObj); + btVector3* pForce = (btVector3*)p_fbuf; + *pForce = rb->getTotalForce(); + p_fbuf[3] = rb->getLinearDamping(); + p_fbuf += 4; + btVector3* pTorque = (btVector3*)p_fbuf; + *pTorque = rb->getTotalTorque(); + p_fbuf[3] = rb->getAngularDamping(); + p_fbuf += 4; + if(m_copyIntegrDataToGPU) + { + for(int k = 0; k < 3; k++) + { + btVector3* pInert = (btVector3*)(p_mbuf + k * 4); + *pInert = rb->getInvInertiaTensorWorld().getRow(k); + } + p_mbuf[3] = rb->getInvMass(); + p_mbuf += 12; + } + btVector3 v = rb->getLinearVelocity(); + m_hVel[i] = *((float4*)&v); + v = rb->getAngularVelocity(); + m_hAngVel[i] = *((float4*)&v); + } + } + if(m_useCPUSolver) + { + //BT_PROFILE("motIntegr -- integrate on CPU"); + btGpu_integrVel(m_hForceTorqueDamp, m_hInvInertiaMass, m_hVel, m_hAngVel, timeStep, m_numObj); + writebackData(); + } + else + { +#ifdef BT_USE_CUDA + //BT_PROFILE("CUDA motIntegr -- integrate on CUDA"); + btCuda_copyArrayToDevice(m_dForceTorqueDamp, m_hForceTorqueDamp, sizeof(float) * m_numObj * 4 * 2); + if(m_copyIntegrDataToGPU) + { + btCuda_copyArrayToDevice(m_dInvInertiaMass, m_hInvInertiaMass, sizeof(float) * m_numObj * 4 * 3); + } + btCuda_copyArrayToDevice(m_dVel, m_hVel, m_numObj * sizeof(float4)); + btCuda_copyArrayToDevice(m_dAngVel, m_hAngVel, m_numObj * sizeof(float4)); + btCuda_integrVel(m_dForceTorqueDamp, m_dInvInertiaMass, m_dVel, m_dAngVel, timeStep, m_numObj); + copyDataFromGPU(); + writebackData(); +#endif + } + } + else + { + btDiscreteDynamicsWorld::predictUnconstraintMotion(timeStep); + m_copyIntegrDataToGPU = true; + } +} // btCudaDemoDynamicsWorld3D::predictUnconstraintMotion() + +//-------------------------------------------------------------------------- + +void btCudaDemoDynamicsWorld3D::integrateTransforms(btScalar timeStep) +{ + if(m_useCudaMotIntegr) + { + BT_PROFILE("motIntegr -- integrateTransforms"); + if(m_useCPUSolver) + { + btGpu_integrTrans(m_hTrans, m_hVel, m_hAngVel, timeStep, m_numObj); + } + else + { +#ifdef BT_USE_CUDA + btCuda_integrTrans(m_dTrans, m_dVel, m_dAngVel, timeStep, m_numObj); + btCuda_copyArrayFromDevice(m_hTrans, m_dTrans, m_numObj * sizeof(float4) * 4); +#endif + } + m_numObj = getNumCollisionObjects(); + for(int i = 0; i < m_numObj; i++) + { + btCollisionObject* colObj = m_collisionObjects[i]; + btRigidBody* rb = btRigidBody::upcast(colObj); + btVector3 v; + btTransform tr; + const btVector3& v0 = *((btVector3*)&m_hTrans[i * 4 + 0]); + const btVector3& v1 = *((btVector3*)&m_hTrans[i * 4 + 1]); + const btVector3& v2 = *((btVector3*)&m_hTrans[i * 4 + 2]); + const btVector3& v3 = *((btVector3*)&m_hTrans[i * 4 + 3]); + tr.getBasis().setValue(v0[0], v1[0], v2[0], v0[1], v1[1], v2[1], v0[2], v1[2], v2[2]); + tr.getOrigin().setValue(v3[0], v3[1], v3[2]); + rb->proceedToTransform(tr); + } + } + else + { + btDiscreteDynamicsWorld::integrateTransforms(timeStep); + } +} // btCudaDemoDynamicsWorld3D::integrateTransforms() + +//-------------------------------------------------------------------------- +//-------------------------------------------------------------------------- +//-------------------------------------------------------------------------- diff --git a/Demos/Gpu3dDemo/btGpuDemoDynamicsWorld3D.h b/Demos/Gpu3dDemo/btGpuDemoDynamicsWorld3D.h new file mode 100644 index 000000000..d13de52bf --- /dev/null +++ b/Demos/Gpu3dDemo/btGpuDemoDynamicsWorld3D.h @@ -0,0 +1,252 @@ +/* +Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org +Copyright (C) 2006, 2007 Sony Computer Entertainment Inc. + +This software is provided 'as-is', without any express or implied warranty. +In no event will the authors be held liable for any damages arising from the use of this software. +Permission is granted to anyone to use this software for any purpose, +including commercial applications, and to alter it and redistribute it freely, +subject to the following restrictions: + +1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. +2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. +3. This notice may not be removed or altered from any source distribution. +*/ + +#ifndef BT_CUDA_DEMO_DYNAMICS_WORLD3D_H +#define BT_CUDA_DEMO_DYNAMICS_WORLD3D_H + +#include "BulletDynamics/Dynamics/btDiscreteDynamicsWorld.h" + +//#define BT_USE_CUDA 1 +// To enable CUDA : +// 1. Uncomment //#define BT_USE_CUDA 1 +// 2. Build and add libbulletcuda (Extras/CUDA) to project +// 3. Add $(CUDA_LIB_PATH) and cudart.lib to linker properties + + +#ifdef BT_USE_CUDA + #include "BulletMultiThreaded/btGpuDefines.h" + #undef BT_GPU_PREF + #define BT_GPU_PREF(func) btCuda_##func + #include "BulletMultiThreaded/btGpuUtilsSharedDefs.h" +#else + #include "BulletMultiThreaded/btGpuDefines.h" + #include "../../src/BulletMultiThreaded/btGpuUtilsSharedDefs.h" +#endif + +#undef BT_GPU_PREF + + +#if 0 // ### +#include +#define BT_GPU_PREF(func) btCuda_##func +#include "../../src/BulletMultiThreaded/btGpuUtilsSharedDefs.h" +#undef BT_GPU_PREF +#endif + +#include "btGpuDemo3DSharedTypes.h" + +//#define CUDA_DEMO_DYNAMICS_WORLD3D_MAX_BATCHES 20 +#define CUDA_DEMO_DYNAMICS_WORLD3D_MAX_BATCHES 15 + +class btCudaDemoDynamicsWorld3D : public btDiscreteDynamicsWorld +{ +protected: + int m_maxObj; + int m_maxNeihbors; + int m_maxConstr; + int m_maxPointsPerConstr; + + int m_numObj; + int m_numSimStep; + bool m_useCPUSolver; + bool m_useSeqImpSolver; + bool m_useCudaMotIntegr; + bool m_copyIntegrDataToGPU; + + +#ifdef BT_USE_CUDA + float4* m_dTrans; + float4* m_dVel; + float4* m_dAngVel; + int2* m_dIds; + int* m_dBatchIds; + float* m_dLambdaDtBox; + float* m_dPositionConstraint; + float3* m_dNormal; + float3* m_dContact; + float* m_dForceTorqueDamp; + float* m_dInvInertiaMass; +#endif + + float4* m_hTrans; + float4* m_hVel; + float4* m_hAngVel; + int* m_hConstraintBuffer; + int* m_hConstraintCounter; + int m_maxBatches; + int m_numBatches; + int m_numConstraints; + int2* m_hIds; + int* m_hBatchIds; + + int m_maxVtxPerObj; + + + // ------------- these are only needed for CPU version and for debugging + float* m_hLambdaDtBox; + float* m_hPositionConstraint; + float3* m_hNormal; + float3* m_hContact; + // ------------- + + btScalar m_objRad; + btVector3 m_worldMin; + btVector3 m_worldMax; + + //------------------------------- + int* m_hConstraintUsed; + + //------------------------------- + + float* m_hForceTorqueDamp; + float* m_hInvInertiaMass; + +public: + int m_numInBatches[CUDA_DEMO_DYNAMICS_WORLD3D_MAX_BATCHES]; + + + btCudaDemoDynamicsWorld3D(btDispatcher* dispatcher,btBroadphaseInterface* pairCache,btConstraintSolver* constraintSolver,btCollisionConfiguration* collisionConfiguration, int maxPointsPerConstr = 4) + : btDiscreteDynamicsWorld(dispatcher, pairCache, constraintSolver, collisionConfiguration) + { + m_useCPUSolver = false; + m_useSeqImpSolver = false; + m_useCudaMotIntegr = true; + m_copyIntegrDataToGPU = true; + m_maxObj = 32768; + m_maxNeihbors = 26; + m_maxConstr = m_maxObj * m_maxNeihbors; + int sz = m_maxConstr; + m_hConstraintBuffer = new int[sz]; + m_hConstraintCounter = new int[m_maxObj]; + m_maxBatches = CUDA_DEMO_DYNAMICS_WORLD3D_MAX_BATCHES; + m_hIds = new int2[sz]; + m_hBatchIds = new int[sz]; + for(int i = 0; i < sz; i++) + { + m_hBatchIds[i] = -1; + } + m_hTrans = new float4[m_maxObj * 4]; + m_hVel = new float4[m_maxObj]; + m_hAngVel = new float4[m_maxObj]; + + m_maxPointsPerConstr = maxPointsPerConstr; + +#ifdef BT_USE_CUDA + btCuda_allocateArray((void**)&m_dTrans, sizeof(float4) * m_maxObj * 4); + btCuda_allocateArray((void**)&m_dVel, sizeof(float4) * m_maxObj); + btCuda_allocateArray((void**)&m_dAngVel, sizeof(float4) * m_maxObj); + + btCuda_allocateArray((void**)&m_dIds, sizeof(int2) * sz); + btCuda_allocateArray((void**)&m_dBatchIds, sizeof(int) * sz); + + + btCuda_allocateArray((void**)&m_dLambdaDtBox, sizeof(float) * sz * m_maxPointsPerConstr); + btCuda_allocateArray((void**)&m_dPositionConstraint, sizeof(float) * sz * m_maxPointsPerConstr); + btCuda_allocateArray((void**)&m_dNormal, sizeof(float3) * sz * m_maxPointsPerConstr); + btCuda_allocateArray((void**)&m_dContact, sizeof(float3) * sz * m_maxPointsPerConstr); + + btCuda_allocateArray((void**)&m_dForceTorqueDamp, sizeof(float) * m_maxObj * 4 * 2); + btCuda_allocateArray((void**)&m_dInvInertiaMass, sizeof(float) * m_maxObj * 4 * 3); +#endif + + m_hLambdaDtBox = new float[sz * m_maxPointsPerConstr]; + m_hPositionConstraint = new float[sz * m_maxPointsPerConstr]; + m_hNormal = new float3[sz * m_maxPointsPerConstr]; + m_hContact = new float3[sz * m_maxPointsPerConstr]; + + m_numSimStep = 0; + + m_objRad = 1.0f; + + m_hConstraintUsed = new int[sz]; + + m_hForceTorqueDamp = new float[m_maxObj * 4 * 2]; + m_hInvInertiaMass = new float[4 * m_maxObj * 3]; + + } + virtual ~btCudaDemoDynamicsWorld3D() + { + delete [] m_hConstraintBuffer; + delete [] m_hConstraintCounter; + delete [] m_hIds; + delete [] m_hBatchIds; + delete [] m_hTrans; + delete [] m_hVel; + delete [] m_hAngVel; + + delete [] m_hLambdaDtBox; + delete [] m_hPositionConstraint; + delete [] m_hNormal; + delete [] m_hContact; + delete [] m_hConstraintUsed; + + delete [] m_hForceTorqueDamp; + delete [] m_hInvInertiaMass; + + +#ifdef BT_USE_CUDA + btCuda_freeArray(m_dTrans); + btCuda_freeArray(m_dVel); + btCuda_freeArray(m_dAngVel); + + btCuda_freeArray(m_dIds); + btCuda_freeArray(m_dBatchIds); + btCuda_freeArray(m_dLambdaDtBox); + btCuda_freeArray(m_dPositionConstraint); + btCuda_freeArray(m_dNormal); + btCuda_freeArray(m_dContact); + btCuda_freeArray(m_dForceTorqueDamp); + btCuda_freeArray(m_dInvInertiaMass); +#endif + + } + virtual void calculateSimulationIslands() + { + if(m_useSeqImpSolver) + { + btDiscreteDynamicsWorld::calculateSimulationIslands(); + } + } + virtual void solveConstraints(btContactSolverInfo& solverInfo); + + virtual void predictUnconstraintMotion(btScalar timeStep); + virtual void integrateTransforms(btScalar timeStep); + + + + + void solveConstraintsCPU(btContactSolverInfo& solverInfo); + + void debugDrawConstraints(int selectedBatch, const float* pColorTab); + + void setObjRad(btScalar rad) { m_objRad = rad; } + void setWorldMin(const btVector3& worldMin) { m_worldMin = worldMin; } + void setWorldMax(const btVector3& worldMax) { m_worldMax = worldMax; } + + void grabData(); + void grabObjData(); + void grabConstrData(); + void createBatches(); + void copyDataToGPU(); + void copyDataFromGPU(); + void writebackData(); + void setUseCPUSolver(bool useCPU) { m_useCPUSolver = useCPU; } + void setUseSeqImpSolver(bool useSeqImpSolver) { m_useSeqImpSolver = useSeqImpSolver; } + void setUseCudaMotIntegr(bool useCudaMotIntegr) { m_useCudaMotIntegr = useCudaMotIntegr; } + void resetScene(void) { m_copyIntegrDataToGPU = true; } +}; + + +#endif //BT_CUDA_DEMO_DYNAMICS_WORLD3D_H diff --git a/Demos/Gpu3dDemo/main.cpp b/Demos/Gpu3dDemo/main.cpp new file mode 100644 index 000000000..b008463ea --- /dev/null +++ b/Demos/Gpu3dDemo/main.cpp @@ -0,0 +1,61 @@ +/* +Bullet Continuous Collision Detection and Physics Library +Copyright (c) 2003-2007 Erwin Coumans http://continuousphysics.com/Bullet/ + +This software is provided 'as-is', without any express or implied warranty. +In no event will the authors be held liable for any damages arising from the use of this software. +Permission is granted to anyone to use this software for any purpose, +including commercial applications, and to alter it and redistribute it freely, +subject to the following restrictions: + +1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. +2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. +3. This notice may not be removed or altered from any source distribution. +*/ + +#include "BasicDemo3D.h" +#include "GlutStuff.h" +#include "GLDebugDrawer.h" +#include "btBulletDynamicsCommon.h" +#include "LinearMath/btHashMap.h" + +class OurValue + { + int m_uid; + + public: + OurValue(const btVector3& initialPos) + :m_position(initialPos) + { + static int gUid=0; + m_uid=gUid; + gUid++; + } + + btVector3 m_position; + int getUid() const + { + return m_uid; + } + }; + + +int main(int argc,char** argv) +{ + GLDebugDrawer gDebugDrawer; + + BasicDemo3D ccdDemo; + ccdDemo.initPhysics(); + ccdDemo.getDynamicsWorld()->setDebugDrawer(&gDebugDrawer); + ccdDemo.setWireMode(false); + + +#ifdef CHECK_MEMORY_LEAKS + ccdDemo.exitPhysics(); +#else + return glutmain(argc, argv,640,480,"Bullet Physics Demo. http://bulletphysics.com",&ccdDemo); +#endif + + //default glut doesn't return from mainloop + return 0; +} \ No newline at end of file diff --git a/Extras/CUDA/btGpuDemo3dCudaFunc.cu b/Extras/CUDA/btGpuDemo3dCudaFunc.cu new file mode 100644 index 000000000..7aec906a0 --- /dev/null +++ b/Extras/CUDA/btGpuDemo3dCudaFunc.cu @@ -0,0 +1,46 @@ +/* +Bullet Continuous Collision Detection and Physics Library, http://bulletphysics.org +Copyright (C) 2006, 2007 Sony Computer Entertainment Inc. + +This software is provided 'as-is', without any express or implied warranty. +In no event will the authors be held liable for any damages arising from the use of this software. +Permission is granted to anyone to use this software for any purpose, +including commercial applications, and to alter it and redistribute it freely, +subject to the following restrictions: + +1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. +2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. +3. This notice may not be removed or altered from any source distribution. +*/ + +#include +#include +#include + +#include "../../Extras/CUDA/cutil_math.h" +#include "math_constants.h" + +#include + +//---------------------------------------------------------------------------------------- + +#include "../../Extras/CUDA/btCudaDefines.h" + +//---------------------------------------------------------------------------------------- + +#include "../../src/BulletMultiThreaded/btGpuUtilsSharedDefs.h" +#include "../../Demos/Gpu3dDemo/btGpuDemo3dSharedTypes.h" +#include "../../Demos/Gpu3dDemo/btGpuDemo3dSharedDefs.h" + +//---------------------------------------------------------------------------------------- + +texture posTex; + +//---------------------------------------------------------------------------------------- + +#include "../../Demos/Gpu3dDemo/btGpuDemo3dSharedCode.h" + +//---------------------------------------------------------------------------------------- +//---------------------------------------------------------------------------------------- +//---------------------------------------------------------------------------------------- + diff --git a/Extras/CUDA/libbulletcuda.vcproj b/Extras/CUDA/libbulletcuda.vcproj index 82518d122..69bd35bef 100644 --- a/Extras/CUDA/libbulletcuda.vcproj +++ b/Extras/CUDA/libbulletcuda.vcproj @@ -584,6 +584,30 @@ /> + + + + + + + +