diff --git a/Demos/DX11ClothDemo/cap.h b/Demos/DX11ClothDemo/cap.h index e2d3d8e81..c4a90b2f2 100644 --- a/Demos/DX11ClothDemo/cap.h +++ b/Demos/DX11ClothDemo/cap.h @@ -13,6 +13,8 @@ subject to the following restrictions: 3. This notice may not be removed or altered from any source distribution. */ +#ifndef CAP_H +#define CAP_H class cap { @@ -35,7 +37,7 @@ class cap void set_collision_object(btCollisionObject* co) { - collisionObject = co; + collisionObject = co; } void set_collision_shape(btCollisionShape* cs) @@ -56,12 +58,20 @@ class cap void destroy() { - + SAFE_RELEASE( g_pIndexBuffer ); + SAFE_RELEASE( pVB[0] ); + SAFE_RELEASE( texture2D ); + SAFE_RELEASE( texture2D_view ); } void draw(void) { + + if (!collisionObject) + return; + + ID3D11DeviceContext* pd3dImmediateContext = DXUTGetD3D11DeviceContext(); D3DXMATRIX mWorldViewProjection; @@ -113,6 +123,7 @@ class cap + btTransform trans = collisionObject->getWorldTransform(); @@ -204,86 +215,84 @@ class cap vertex_struct *vertices = new vertex_struct[width*height]; btCapsuleShape* cs = static_cast(collisionShape); - float radius = cs->getRadius(); - float halfHeight = cs->getHalfHeight(); - - - if (top) + if (cs) { - for(int y = 0; y < height; y++) - { - for(int x = 0; x < width; x++) - { - float X = (x/((float)(width-1)))*3.14159; - float Y = (y/((float)(height-1)))*3.14159; - float z_coord = radius*cos(X)*sin(Y); - float y_coord = radius*sin(X)*sin(Y) + halfHeight; - float x_coord = radius*cos(Y); - vertices[y*width+x].Pos = D3DXVECTOR3(x_coord, y_coord, z_coord); - vertices[y*width+x].Normal = D3DXVECTOR3(x_coord,y_coord-halfHeight,z_coord); - vertices[y*width+x].Texcoord = D3DXVECTOR2(x/( (float)(width-1)), y/((float)(height-1))); - } - } + float radius = cs->getRadius(); + float halfHeight = cs->getHalfHeight(); + - } - - else - - { - - for(int y = 0; y < height; y++) - { - for(int x = 0; x < width; x++) - { - float X = (x/((float)(width-1)))*3.14159; - float Y = (y/((float)(height-1)))*3.14159; - float z_coord = radius*cos(X)*sin(Y); - float y_coord = -radius*sin(X)*sin(Y) - halfHeight; - float x_coord = radius*cos(Y); - vertices[y*width+x].Pos = D3DXVECTOR3(x_coord, y_coord, z_coord); - vertices[y*width+x].Normal = D3DXVECTOR3(x_coord,y_coord+halfHeight,z_coord); - vertices[y*width+x].Texcoord = D3DXVECTOR2(x/( (float)(width-1)), y/((float)(height-1))); - } - } - - } - - D3D11_SUBRESOURCE_DATA InitData; - InitData.pSysMem = vertices; - InitData.SysMemPitch = 0; - InitData.SysMemSlicePitch = 0; - - HRESULT hr = g_pd3dDevice->CreateBuffer(&bufferDesc, &InitData, &pVB[0]); - - - - //What is this vertex stride thing all about? - Strides[0] = ( UINT )g_Mesh11.GetVertexStride( 0, 0 ); - Offsets[0] = 0; - - unsigned int* indices = new unsigned int[width*3*2+2 + height*width*3*2]; - - for(int y = 0; y < height-1; y++) - { - for(int x = 0; x < width-1; x++) + if (top) { - indices[x*3*2 + y*width*3*2] = x + y*width; - indices[x*3*2+1 + y*width*3*2] = x+1 + y*width; - indices[x*3*2+2 + y*width*3*2] = x+width + y*width; - - indices[x*3*2 + 3 + y*width*3*2] = x + 1 + y*width; - indices[x*3*2 + 4 + y*width*3*2] = x+(width+1) + y*width; - indices[x*3*2 + 5 + y*width*3*2] = x+width + y*width; - + for(int y = 0; y < height; y++) + { + for(int x = 0; x < width; x++) + { + float X = (x/((float)(width-1)))*3.14159; + float Y = (y/((float)(height-1)))*3.14159; + float z_coord = radius*cos(X)*sin(Y); + float y_coord = radius*sin(X)*sin(Y) + halfHeight; + float x_coord = radius*cos(Y); + vertices[y*width+x].Pos = D3DXVECTOR3(x_coord, y_coord, z_coord); + vertices[y*width+x].Normal = D3DXVECTOR3(x_coord,y_coord-halfHeight,z_coord); + vertices[y*width+x].Texcoord = D3DXVECTOR2(x/( (float)(width-1)), y/((float)(height-1))); + } + } + } else { + for(int y = 0; y < height; y++) + { + for(int x = 0; x < width; x++) + { + float X = (x/((float)(width-1)))*3.14159; + float Y = (y/((float)(height-1)))*3.14159; + float z_coord = radius*cos(X)*sin(Y); + float y_coord = -radius*sin(X)*sin(Y) - halfHeight; + float x_coord = radius*cos(Y); + vertices[y*width+x].Pos = D3DXVECTOR3(x_coord, y_coord, z_coord); + vertices[y*width+x].Normal = D3DXVECTOR3(x_coord,y_coord+halfHeight,z_coord); + vertices[y*width+x].Texcoord = D3DXVECTOR2(x/( (float)(width-1)), y/((float)(height-1))); + } + } } + + D3D11_SUBRESOURCE_DATA InitData; + InitData.pSysMem = vertices; + InitData.SysMemPitch = 0; + InitData.SysMemSlicePitch = 0; + + HRESULT hr = g_pd3dDevice->CreateBuffer(&bufferDesc, &InitData, &pVB[0]); + + + + //What is this vertex stride thing all about? + Strides[0] = ( UINT )g_Mesh11.GetVertexStride( 0, 0 ); + Offsets[0] = 0; + + unsigned int* indices = new unsigned int[width*3*2+2 + height*width*3*2]; + + for(int y = 0; y < height-1; y++) + { + for(int x = 0; x < width-1; x++) + { + indices[x*3*2 + y*width*3*2] = x + y*width; + indices[x*3*2+1 + y*width*3*2] = x+1 + y*width; + indices[x*3*2+2 + y*width*3*2] = x+width + y*width; + + indices[x*3*2 + 3 + y*width*3*2] = x + 1 + y*width; + indices[x*3*2 + 4 + y*width*3*2] = x+(width+1) + y*width; + indices[x*3*2 + 5 + y*width*3*2] = x+width + y*width; + + } + } + + bufferDesc.ByteWidth = sizeof(unsigned int)*(width*3*2+2 + height*width*3*2); + bufferDesc.BindFlags = D3D11_BIND_INDEX_BUFFER; + + InitData.pSysMem = indices; + + hr = g_pd3dDevice->CreateBuffer(&bufferDesc, &InitData, &g_pIndexBuffer); + hr = hr; } - - bufferDesc.ByteWidth = sizeof(unsigned int)*(width*3*2+2 + height*width*3*2); - bufferDesc.BindFlags = D3D11_BIND_INDEX_BUFFER; - - InitData.pSysMem = indices; - - hr = g_pd3dDevice->CreateBuffer(&bufferDesc, &InitData, &g_pIndexBuffer); - hr = hr; } }; + +#endif CAP_H diff --git a/Demos/DX11ClothDemo/cloth.h b/Demos/DX11ClothDemo/cloth.h index c130548c1..28a329b54 100644 --- a/Demos/DX11ClothDemo/cloth.h +++ b/Demos/DX11ClothDemo/cloth.h @@ -13,6 +13,8 @@ subject to the following restrictions: 3. This notice may not be removed or altered from any source distribution. */ +#ifndef CLOTH_H +#define CLOTH_H #include #include @@ -304,3 +306,6 @@ public: } }; +#endif //CLOTH_H + + diff --git a/Demos/DX11ClothDemo/cloth_renderer.cpp b/Demos/DX11ClothDemo/cloth_renderer.cpp index be86c41d3..2c05b00b4 100644 --- a/Demos/DX11ClothDemo/cloth_renderer.cpp +++ b/Demos/DX11ClothDemo/cloth_renderer.cpp @@ -38,15 +38,18 @@ class btDX11SIMDAwareSoftBodySolver; #include "BulletSoftBody/btSoftBodyRigidBodyCollisionConfiguration.h" #define USE_SIMDAWARE_SOLVER -//#define USE_GPU_SOLVER +#define USE_GPU_SOLVER #define USE_GPU_COPY const int numFlags = 5; const int clothWidth = 40; -const int clothHeight = 60;//60; +const int clothHeight = 60; float _windAngle = 1.0;//0.4; float _windStrength = 15; +//#define TABLETEST + + #include #include @@ -171,21 +174,20 @@ struct vertex_struct +#include "capsule.h" #include "cap.h" #include "cylinder.h" #include "cloth.h" -//#include "capsule.h" //cylinder cyl_1; //cap cap_1; btRigidBody *capCollider; - +capsule my_capsule; btAlignedObjectArray cloths; -//capsule my_capsule; ////////////////////////////////////////// // Bullet globals @@ -205,6 +207,8 @@ btCPUSoftBodySolver *g_cpuSolver = NULL; btDX11SoftBodySolver *g_dx11Solver = NULL; btDX11SIMDAwareSoftBodySolver *g_dx11SIMDSolver = NULL; +btSoftBodySolverOutput *g_softBodyOutput = NULL; + btSoftBodySolver *g_solver = NULL; // End bullet globals @@ -359,13 +363,23 @@ void createFlag( int width, int height, btAlignedObjectArray &flag defaultRotate[0] = btVector3(cos(rotateAngleRoundZ), sin(rotateAngleRoundZ), 0.f); defaultRotate[1] = btVector3(-sin(rotateAngleRoundZ), cos(rotateAngleRoundZ), 0.f); defaultRotate[2] = btVector3(0.f, 0.f, 1.f); + + + //btMatrix3x3 defaultRotateAndScale( (defaultRotateX*defaultRotate) ); +#ifdef TABLETEST + btMatrix3x3 defaultRotateX; + rotateAngleRoundX = 3.141592654/2; + defaultRotateX[0] = btVector3(1.f, 0.f, 0.f); + defaultRotateX[1] = btVector3( 0.f, cos(rotateAngleRoundX), sin(rotateAngleRoundX)); + defaultRotateX[2] = btVector3(0.f, -sin(rotateAngleRoundX), cos(rotateAngleRoundX)); + btMatrix3x3 defaultRotateAndScale( (defaultRotateX) ); +#else btMatrix3x3 defaultRotateX; defaultRotateX[0] = btVector3(1.f, 0.f, 0.f); defaultRotateX[1] = btVector3( 0.f, cos(rotateAngleRoundX), sin(rotateAngleRoundX)); defaultRotateX[2] = btVector3(0.f, -sin(rotateAngleRoundX), cos(rotateAngleRoundX)); - - //btMatrix3x3 defaultRotateAndScale( (defaultRotateX*defaultRotate) ); btMatrix3x3 defaultRotateAndScale( (defaultRotateX) ); +#endif // Construct the sequence flags applying a slightly different translation to each one to arrange them @@ -387,10 +401,12 @@ void createFlag( int width, int height, btAlignedObjectArray &flag softBody->setMass(i, 10.f/mesh.m_numVertices); } +#ifndef TABLETEST // Set the fixed points softBody->setMass((height-1)*(width), 0.f); softBody->setMass((height-1)*(width) + width - 1, 0.f); softBody->setMass((height-1)*width + width/2, 0.f); +#endif softBody->m_cfg.collisions = btSoftBody::fCollision::CL_SS+btSoftBody::fCollision::CL_RS; softBody->m_cfg.kLF = 0.0005f; @@ -401,7 +417,7 @@ void createFlag( int width, int height, btAlignedObjectArray &flag flags.push_back( softBody ); softBody->transform( transform ); - + softBody->setFriction( 0.8f ); m_dynamicsWorld->addSoftBody( softBody ); } @@ -437,11 +453,15 @@ void updatePhysicsWorld() cloth->setWindVelocity( btVector3(xCoordinate, 0, zCoordinate) ); } } +#ifndef TABLETEST + if (capCollider) + { + btVector3 origin( capCollider->getWorldTransform().getOrigin() ); + origin.setZ( origin.getZ() + 0.01 ); + capCollider->getWorldTransform().setOrigin( origin ); + } +#endif - //btVector3 origin( capCollider->getWorldTransform().getOrigin() ); - //origin.setX( origin.getX() + 0.05 ); - //capCollider->getWorldTransform().setOrigin( origin ); - counter++; } @@ -452,13 +472,25 @@ void initBullet(void) #ifdef USE_GPU_SOLVER g_dx11Solver = new btDX11SoftBodySolver( g_pd3dDevice, DXUTGetD3D11DeviceContext() ); g_solver = g_dx11Solver; +#ifdef USE_GPU_COPY + g_softBodyOutput = new btSoftBodySolverOutputDXtoDX( g_pd3dDevice, DXUTGetD3D11DeviceContext() ); +#else // #ifdef USE_GPU_COPY + g_softBodyOutput = new btSoftBodySolverOutputDXtoCPU; +#endif // #ifdef USE_GPU_COPY #else #ifdef USE_SIMDAWARE_SOLVER g_dx11SIMDSolver = new btDX11SIMDAwareSoftBodySolver( g_pd3dDevice, DXUTGetD3D11DeviceContext() ); g_solver = g_dx11SIMDSolver; + g_softBodyOutput = new btSoftBodySolverOutputDXtoCPU; +#ifdef USE_GPU_COPY + g_softBodyOutput = new btSoftBodySolverOutputDXtoDX( g_pd3dDevice, DXUTGetD3D11DeviceContext() ); +#else // #ifdef USE_GPU_COPY + g_softBodyOutput = new btSoftBodySolverOutputDXtoCPU; +#endif // #ifdef USE_GPU_COPY #else g_cpuSolver = new btCPUSoftBodySolver; g_solver = g_cpuSolver; + g_softBodyOutput = new btSoftBodySolverOutputCPUtoCPU; //g_defaultSolver = new btDefaultSoftBodySolver; //g_solver = g_defaultSolver; #endif @@ -519,7 +551,8 @@ void initBullet(void) //rigidbody is dynamic if and only if mass is non zero, otherwise static bool isDynamic = (mass != 0.f); - btCollisionShape *capsuleShape = new btCapsuleShape(5, 30); + btCollisionShape *capsuleShape = new btCapsuleShape(5, 10); + capsuleShape->setMargin( 0.5 ); @@ -532,11 +565,21 @@ void initBullet(void) m_collisionShapes.push_back(capsuleShape); btTransform capsuleTransform; capsuleTransform.setIdentity(); - capsuleTransform.setOrigin(btVector3(0, 10, 0)); +#ifdef TABLETEST + capsuleTransform.setOrigin(btVector3(0, 10, -11)); + const btScalar pi = 3.141592654; + capsuleTransform.setRotation(btQuaternion(0, 0, pi/2)); +#else + capsuleTransform.setOrigin(btVector3(0, 20, -10)); + + const btScalar pi = 3.141592654; + //capsuleTransform.setRotation(btQuaternion(0, 0, pi/2)); + capsuleTransform.setRotation(btQuaternion(0, 0, 0)); +#endif btDefaultMotionState* myMotionState = new btDefaultMotionState(capsuleTransform); btRigidBody::btRigidBodyConstructionInfo rbInfo(mass,myMotionState,capsuleShape,localInertia); btRigidBody* body = new btRigidBody(rbInfo); - + body->setFriction( 0.8f ); my_capsule.set_collision_object(body); m_dynamicsWorld->addRigidBody(body); @@ -644,6 +687,7 @@ int WINAPI wWinMain( HINSTANCE hInstance, HINSTANCE hPrevInstance, LPWSTR lpCmdL // Enable run-time memory check for debug builds. #if defined(DEBUG) | defined(_DEBUG) _CrtSetDbgFlag( _CRTDBG_ALLOC_MEM_DF | _CRTDBG_LEAK_CHECK_DF ); + _CrtSetReportMode ( _CRT_ERROR, _CRTDBG_MODE_DEBUG); #endif // DXUT will create and use the best device (either D3D9 or D3D11) @@ -694,10 +738,10 @@ void InitApp() g_SampleUI.Init( &g_DialogResourceManager ); g_HUD.SetCallback( OnGUIEvent ); int iY = 10; -//restart is broken, @todo: fix -// g_HUD.AddButton( IDC_TOGGLEFULLSCREEN, L"Toggle full screen", 0, iY, 170, 23 ); -// g_HUD.AddButton( IDC_TOGGLEREF, L"Toggle REF (F3)", 0, iY += 26, 170, 23, VK_F3 ); -// g_HUD.AddButton( IDC_CHANGEDEVICE, L"Change device (F2)", 0, iY += 26, 170, 23, VK_F2 ); + g_HUD.AddButton( IDC_TOGGLEFULLSCREEN, L"Toggle full screen", 0, iY, 170, 23 ); + + g_HUD.AddButton( IDC_TOGGLEREF, L"Toggle REF (F3)", 0, iY += 26, 170, 23, VK_F3 ); + g_HUD.AddButton( IDC_CHANGEDEVICE, L"Change device (F2)", 0, iY += 26, 170, 23, VK_F2 ); g_HUD.AddButton( IDC_PAUSE, L"Pause", 0, iY += 26, 170, 23 ); g_HUD.AddButton( IDC_WIREFRAME, L"Wire frame", 0, iY += 26, 170, 23 ); @@ -1020,8 +1064,8 @@ HRESULT CALLBACK OnD3D11CreateDevice( ID3D11Device* pd3dDevice, const DXGI_SURFA // Setup the camera's view parameters - D3DXVECTOR3 vecEye( 30.0f, -10.0f, -80.0f ); - D3DXVECTOR3 vecAt ( 0.0f, 20.0f, -0.0f ); + D3DXVECTOR3 vecEye( 0.0f, 0.0f, -100.0f ); + D3DXVECTOR3 vecAt ( 0.0f, 0.0f, -0.0f ); g_Camera.SetViewParams( &vecEye, &vecAt ); @@ -1035,25 +1079,12 @@ HRESULT CALLBACK OnD3D11CreateDevice( ID3D11Device* pd3dDevice, const DXGI_SURFA initBullet(); - - //my_capsule.create_buffers(50,40); - - - - std::wstring flagTexsName[] = { - L"atiFlag.bmp", + std::wstring flagTexs[] = { L"amdFlag.bmp", + L"atiFlag.bmp", }; int numFlagTexs = 2; - - - WCHAR flagTexs[2][MAX_PATH]; - - HRESULT res = DXUTFindDXSDKMediaFileCch(flagTexs[0],MAX_PATH, flagTexsName[0].c_str()); - res = DXUTFindDXSDKMediaFileCch(flagTexs[1],MAX_PATH, flagTexsName[1].c_str()); - - for( int flagIndex = 0; flagIndex < numFlags; ++flagIndex ) { cloths[flagIndex].create_texture(flagTexs[flagIndex % numFlagTexs]); @@ -1062,12 +1093,10 @@ HRESULT CALLBACK OnD3D11CreateDevice( ID3D11Device* pd3dDevice, const DXGI_SURFA cloths[flagIndex].z_offset = 0; } - //cap_1.create_texture(); - //cap_1.x_offset = 0; - //cap_1.y_offset = 0; - //cap_1.z_offset = 0; + - //my_capsule.create_texture(); + my_capsule.create_buffers(50,40); + my_capsule.create_texture(); //Turn off backface culling D3D11_RASTERIZER_DESC rsDesc; @@ -1142,10 +1171,7 @@ HRESULT CALLBACK OnD3D11ResizedSwapChain( ID3D11Device* pd3dDevice, IDXGISwapCha } -#ifndef BT_NO_PROFILE btClock m_clock; -#endif //BT_NO_PROFILE - //-------------------------------------------------------------------------------------- // Render the scene using the D3D11 device //-------------------------------------------------------------------------------------- @@ -1157,12 +1183,8 @@ void CALLBACK OnD3D11FrameRender( ID3D11Device* pd3dDevice, ID3D11DeviceContext* //float ms = getDeltaTimeMicroseconds(); -#ifndef BT_NO_PROFILE btScalar dt = (btScalar)m_clock.getTimeMicroseconds(); m_clock.reset(); -#else - btScalar dt = 1000000.f/60.f; -#endif //BT_NO_PROFILE ///step the simulation if (m_dynamicsWorld && !paused) @@ -1198,12 +1220,11 @@ void CALLBACK OnD3D11FrameRender( ID3D11Device* pd3dDevice, ID3D11DeviceContext* for( int flagIndex = 0; flagIndex < m_flags.size(); ++flagIndex ) { - g_solver->copySoftBodyToVertexBuffer( m_flags[flagIndex], cloths[flagIndex].m_vertexBufferDescriptor ); + g_softBodyOutput->copySoftBodyToVertexBuffer( m_flags[flagIndex], cloths[flagIndex].m_vertexBufferDescriptor ); cloths[flagIndex].draw(); } - //my_capsule.draw(); - //cap_1.draw(); + my_capsule.draw(); DXUT_BeginPerfEvent( DXUT_PERFEVENTCOLOR, L"HUD / Stats" ); @@ -1265,6 +1286,8 @@ void CALLBACK OnD3D11DestroyDevice( void* pUserContext ) cloths[flagIndex].destroy(); } + my_capsule.destroy(); + // Shouldn't need to delete this as it's just a soft body and will be deleted later by the collision object cleanup. //for( int flagIndex = 0; flagIndex < m_flags.size(); ++flagIndex ) //{ @@ -1280,6 +1303,8 @@ void CALLBACK OnD3D11DestroyDevice( void* pUserContext ) delete g_dx11Solver; if( g_dx11SIMDSolver ) delete g_dx11SIMDSolver; + if( g_softBodyOutput ) + delete g_softBodyOutput; for(int i=0; i< m_collisionShapes.size(); i++) diff --git a/Demos/DX11ClothDemo/cylinder.h b/Demos/DX11ClothDemo/cylinder.h index 517fcdf51..225cb204a 100644 --- a/Demos/DX11ClothDemo/cylinder.h +++ b/Demos/DX11ClothDemo/cylinder.h @@ -13,6 +13,8 @@ subject to the following restrictions: 3. This notice may not be removed or altered from any source distribution. */ +#ifndef CYLINDER_H +#define CYLINDER_H class cylinder { @@ -48,7 +50,15 @@ class cylinder collisionShape = cs; } + + void destroy() + { + SAFE_RELEASE( g_pIndexBuffer ); + SAFE_RELEASE( pVB[0] ); + SAFE_RELEASE( texture2D ); + SAFE_RELEASE( texture2D_view ); + } void create_texture(void) { @@ -58,13 +68,14 @@ class cylinder loadInfo.Format = DXGI_FORMAT_BC1_UNORM; HRESULT hr = D3DX11CreateShaderResourceViewFromFile(g_pd3dDevice, L"texture.bmp", &loadInfo, NULL, &texture2D_view, NULL); - hr = hr; - } void draw(void) { + if (!collisionObject) + return; + ID3D11DeviceContext* pd3dImmediateContext = DXUTGetD3D11DeviceContext(); D3DXMATRIX mWorldViewProjection; @@ -204,79 +215,83 @@ class cylinder vertex_struct *vertices = new vertex_struct[width*height]; btCapsuleShape* cs = static_cast(collisionShape); - float radius = cs->getRadius(); - float halfHeight = cs->getHalfHeight(); - - for(int y = 0; y < height; y++) + if (cs) { - for(int x = 0; x < width; x++) + float radius = cs->getRadius(); + float halfHeight = cs->getHalfHeight(); + + for(int y = 0; y < height; y++) { - double coord_2 = sin(2.2*3.141159*y/(float)height)*radius; - double coord_1 = cos(2.2*3.141159*y/(float)height)*radius; - //double coord_2 = (y/((float)(height-1)))*1000; - - //coord = sin(y/); + for(int x = 0; x < width; x++) + { + double coord_2 = sin(2.2*3.141159*y/(float)height)*radius; + double coord_1 = cos(2.2*3.141159*y/(float)height)*radius; + //double coord_2 = (y/((float)(height-1)))*1000; + + //coord = sin(y/); - vertices[y*width+x].Pos = D3DXVECTOR3(coord_1, ((x/((float)(width-1)))-.5)*halfHeight*2, coord_2); - vertices[y*width+x].Normal = D3DXVECTOR3(coord_1,0,coord_2); - vertices[y*width+x].Texcoord = D3DXVECTOR2(x/( (float)(width-1)), y/((float)(height-1))); + vertices[y*width+x].Pos = D3DXVECTOR3(coord_1, ((x/((float)(width-1)))-.5)*halfHeight*2, coord_2); + vertices[y*width+x].Normal = D3DXVECTOR3(coord_1,0,coord_2); + vertices[y*width+x].Texcoord = D3DXVECTOR2(x/( (float)(width-1)), y/((float)(height-1))); + } } - } - + - /* - for(int y = 0; y < height; y++) - { - for(int x = 0; x < width; x++) + /* + for(int y = 0; y < height; y++) { - double coord = sin(x/5.0)*50; - //coord = sin(y/); + for(int x = 0; x < width; x++) + { + double coord = sin(x/5.0)*50; + //coord = sin(y/); - vertices[y*width+x].Pos = D3DXVECTOR3( (x/((float)(width-1)))*1000, coord, (y/((float)(height-1)))*1000); - vertices[y*width+x].Normal = D3DXVECTOR3(1,0,0); - vertices[y*width+x].Texcoord = D3DXVECTOR2(x/( (float)(width-1)), y/((float)(height-1))); + vertices[y*width+x].Pos = D3DXVECTOR3( (x/((float)(width-1)))*1000, coord, (y/((float)(height-1)))*1000); + vertices[y*width+x].Normal = D3DXVECTOR3(1,0,0); + vertices[y*width+x].Texcoord = D3DXVECTOR2(x/( (float)(width-1)), y/((float)(height-1))); + } } - } - */ + */ - D3D11_SUBRESOURCE_DATA InitData; - InitData.pSysMem = vertices; - InitData.SysMemPitch = 0; - InitData.SysMemSlicePitch = 0; + D3D11_SUBRESOURCE_DATA InitData; + InitData.pSysMem = vertices; + InitData.SysMemPitch = 0; + InitData.SysMemSlicePitch = 0; - HRESULT hr = g_pd3dDevice->CreateBuffer(&bufferDesc, &InitData, &pVB[0]); - - - //What is this vertex stride thing all about? - Strides[0] = ( UINT )g_Mesh11.GetVertexStride( 0, 0 ); - Offsets[0] = 0; + HRESULT hr = g_pd3dDevice->CreateBuffer(&bufferDesc, &InitData, &pVB[0]); + + + //What is this vertex stride thing all about? + Strides[0] = ( UINT )g_Mesh11.GetVertexStride( 0, 0 ); + Offsets[0] = 0; - //unsigned int indices[] = {0,1,2, 1,3,2}; - unsigned int* indices = new unsigned int[width*3*2+2 + height*width*3*2]; + //unsigned int indices[] = {0,1,2, 1,3,2}; + unsigned int* indices = new unsigned int[width*3*2+2 + height*width*3*2]; - for(int y = 0; y < height-1; y++) - { - for(int x = 0; x < width-1; x++) + for(int y = 0; y < height-1; y++) { - indices[x*3*2 + y*width*3*2] = x + y*width; - indices[x*3*2+1 + y*width*3*2] = x+1 + y*width; - indices[x*3*2+2 + y*width*3*2] = x+width + y*width; + for(int x = 0; x < width-1; x++) + { + indices[x*3*2 + y*width*3*2] = x + y*width; + indices[x*3*2+1 + y*width*3*2] = x+1 + y*width; + indices[x*3*2+2 + y*width*3*2] = x+width + y*width; - indices[x*3*2 + 3 + y*width*3*2] = x + 1 + y*width; - indices[x*3*2 + 4 + y*width*3*2] = x+(width+1) + y*width; - indices[x*3*2 + 5 + y*width*3*2] = x+width + y*width; + indices[x*3*2 + 3 + y*width*3*2] = x + 1 + y*width; + indices[x*3*2 + 4 + y*width*3*2] = x+(width+1) + y*width; + indices[x*3*2 + 5 + y*width*3*2] = x+width + y*width; + } } + + bufferDesc.ByteWidth = sizeof(unsigned int)*(width*3*2+2 + height*width*3*2); + bufferDesc.BindFlags = D3D11_BIND_INDEX_BUFFER; + + InitData.pSysMem = indices; + + hr = g_pd3dDevice->CreateBuffer(&bufferDesc, &InitData, &g_pIndexBuffer); + hr = hr; } - - bufferDesc.ByteWidth = sizeof(unsigned int)*(width*3*2+2 + height*width*3*2); - bufferDesc.BindFlags = D3D11_BIND_INDEX_BUFFER; - - InitData.pSysMem = indices; - - hr = g_pd3dDevice->CreateBuffer(&bufferDesc, &InitData, &g_pIndexBuffer); - hr = hr; } }; +#endif CYLINDER_H diff --git a/Demos/OpenCLClothDemo/AMD/CMakeLists.txt b/Demos/OpenCLClothDemo/AMD/CMakeLists.txt index 599cccccc..ad6069953 100644 --- a/Demos/OpenCLClothDemo/AMD/CMakeLists.txt +++ b/Demos/OpenCLClothDemo/AMD/CMakeLists.txt @@ -42,12 +42,14 @@ IF (USE_GLUT) ${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOclCommon.h ${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOclUtils.cpp ${BULLET_PHYSICS_SOURCE_DIR}/Demos/SharedOpenCL/btOclCommon.cpp +# ${BULLET_PHYSICS_SOURCE_DIR}/Demos/OpenGL/GLDebugDrawer.cpp ../gl_win.cpp ../clstuff.cpp ../bmpLoader.cpp ../bmpLoader.h ../clstuff.h ../gl_win.h + ../cloth.h ) ELSE (USE_GLUT) diff --git a/Demos/OpenCLClothDemo/cl_cloth_demo.cpp b/Demos/OpenCLClothDemo/cl_cloth_demo.cpp index 786fa810d..2afa2b2fd 100644 --- a/Demos/OpenCLClothDemo/cl_cloth_demo.cpp +++ b/Demos/OpenCLClothDemo/cl_cloth_demo.cpp @@ -17,12 +17,22 @@ subject to the following restrictions: #include #endif +#ifndef USE_MINICL +#define USE_SIMDAWARE_SOLVER +#define USE_GPU_SOLVER +#define USE_GPU_COPY +#endif //USE_MINICL + + + + #include "clstuff.h" #include "gl_win.h" #include "cloth.h" -#define USE_GPU_SOLVER +#include "../OpenGL/GLDebugDrawer.h" +GLDebugDrawer debugDraw; const int numFlags = 5; const int clothWidth = 40; @@ -32,8 +42,8 @@ float _windStrength = 15; -#include -using namespace std; + + @@ -44,6 +54,13 @@ using namespace std; #include "vectormath/vmInclude.h" #include "BulletMultiThreaded/GpuSoftBodySolvers/CPU/btSoftBodySolver_CPU.h" #include "BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCL.h" +#include "BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCLSIMDAware.h" +#include "BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolverVertexBuffer_OpenGL.h" +#include "BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolverOutputCLtoGL.h" + + +btRigidBody *capCollider; + using Vectormath::Aos::Vector3; @@ -56,6 +73,8 @@ class btConstraintSolver; struct btCollisionAlgorithmCreateFunc; class btDefaultCollisionConfiguration; +#include "BulletSoftBody/btSoftBodyRigidBodyCollisionConfiguration.h" + namespace Vectormath { namespace Aos @@ -73,16 +92,19 @@ btDefaultCollisionConfiguration* m_collisionConfiguration; btCPUSoftBodySolver *g_cpuSolver = NULL; btOpenCLSoftBodySolver *g_openCLSolver = NULL; +btOpenCLSoftBodySolverSIMDAware *g_openCLSIMDSolver = NULL; btSoftBodySolver *g_solver = NULL; +btSoftBodySolverOutput *g_softBodyOutput = NULL; + btAlignedObjectArray m_flags; btSoftRigidDynamicsWorld* m_dynamicsWorld; btAlignedObjectArray cloths; extern cl_context g_cxMainContext; -extern cl_device_id g_cdDevice; -extern cl_command_queue g_cqCommandQue; +extern cl_device_id g_cdDevice; +extern cl_command_queue g_cqCommandQue; const float flagSpacing = 30.f; @@ -254,6 +276,7 @@ void createFlag( btSoftBodySolver &solver, int width, int height, btAlignedObjec btVector3 defaultTranslate(0.f, 20.f, zTranslate); btTransform transform( defaultRotateAndScale, defaultTranslate ); + transform.setOrigin(defaultTranslate); btSoftBody *softBody = createFromIndexedMesh( vertexArray, mesh.m_numVertices, triangleVertexIndexArray, mesh.m_numTriangles, true ); @@ -268,6 +291,11 @@ void createFlag( btSoftBodySolver &solver, int width, int height, btAlignedObjec softBody->setMass((height-1)*width + width/2, 0.f); softBody->m_cfg.collisions = btSoftBody::fCollision::CL_SS+btSoftBody::fCollision::CL_RS; + softBody->m_cfg.kLF = 0.0005f; + softBody->m_cfg.kVCF = 0.001f; + softBody->m_cfg.kDP = 0.f; + softBody->m_cfg.kDG = 0.f; + flags.push_back( softBody ); @@ -283,7 +311,7 @@ void createFlag( btSoftBodySolver &solver, int width, int height, btAlignedObjec void updatePhysicsWorld() { - static int counter = 0; + static int counter = 1; // Change wind velocity a bit based on a frame counter if( (counter % 400) == 0 ) @@ -317,16 +345,33 @@ void initBullet(void) { #ifdef USE_GPU_SOLVER - g_openCLSolver = new btOpenCLSoftBodySolver( g_cqCommandQue, g_cxMainContext); - //g_openCLSolver->setDefaultWorkgroupSize(32); - +#ifdef USE_SIMDAWARE_SOLVER + g_openCLSIMDSolver = new btOpenCLSoftBodySolverSIMDAware( g_cqCommandQue, g_cxMainContext); + g_solver = g_openCLSIMDSolver; +#ifdef USE_GPU_COPY + g_softBodyOutput = new btSoftBodySolverOutputCLtoGL(g_cqCommandQue, g_cxMainContext); +#else // #ifdef USE_GPU_COPY + g_softBodyOutput = new btSoftBodySolverOutputCLtoCPU; +#endif // #ifdef USE_GPU_COPY +#else + g_openCLSolver = new btOpenCLSoftBodySolver( g_cqCommandQue, g_cxMainContext ); g_solver = g_openCLSolver; +#ifdef USE_GPU_COPY + g_softBodyOutput = new btSoftBodySolverOutputCLtoGL(g_cqCommandQue, g_cxMainContext); +#else // #ifdef USE_GPU_COPY + g_softBodyOutput = new btSoftBodySolverOutputCLtoCPU; +#endif // #ifdef USE_GPU_COPY +#endif #else g_cpuSolver = new btCPUSoftBodySolver; g_solver = g_cpuSolver; + g_softBodyOutput = new btSoftBodySolverOutputCPUtoCPU; #endif - m_collisionConfiguration = new btDefaultCollisionConfiguration(); + //m_collisionConfiguration = new btDefaultCollisionConfiguration(); + m_collisionConfiguration = new btSoftBodyRigidBodyCollisionConfiguration(); + + m_dispatcher = new btCollisionDispatcher(m_collisionConfiguration); m_broadphase = new btDbvtBroadphase(); btSequentialImpulseConstraintSolver* sol = new btSequentialImpulseConstraintSolver; @@ -376,6 +421,51 @@ void initBullet(void) #endif + #if 1 + { + btScalar mass(0.); + + //btScalar mass(1.); + + //rigidbody is dynamic if and only if mass is non zero, otherwise static + bool isDynamic = (mass != 0.f); + + btCollisionShape *capsuleShape = new btCapsuleShape(5, 10); + capsuleShape->setMargin( 0.5 ); + + + + + btVector3 localInertia(0,0,0); + if (isDynamic) + capsuleShape->calculateLocalInertia(mass,localInertia); + + m_collisionShapes.push_back(capsuleShape); + btTransform capsuleTransform; + capsuleTransform.setIdentity(); +#ifdef TABLETEST + capsuleTransform.setOrigin(btVector3(0, 10, -11)); + const btScalar pi = 3.141592654; + capsuleTransform.setRotation(btQuaternion(0, 0, pi/2)); +#else + capsuleTransform.setOrigin(btVector3(0, 0, 0)); + + const btScalar pi = 3.141592654; + //capsuleTransform.setRotation(btQuaternion(0, 0, pi/2)); + capsuleTransform.setRotation(btQuaternion(0, 0, 0)); +#endif + btDefaultMotionState* myMotionState = new btDefaultMotionState(capsuleTransform); + btRigidBody::btRigidBodyConstructionInfo rbInfo(mass,myMotionState,capsuleShape,localInertia); + btRigidBody* body = new btRigidBody(rbInfo); + body->setFriction( 0.8f ); + + m_dynamicsWorld->addRigidBody(body); + //cap_1.collisionShape = body; + capCollider = body; + } +#endif + + #ifdef USE_GPU_SOLVER createFlag( *g_openCLSolver, clothWidth, clothHeight, m_flags ); #else @@ -391,7 +481,12 @@ void initBullet(void) // In this case we have a DX11 output buffer with a vertex at index 0, 8, 16 and so on as well as a normal at 3, 11, 19 etc. // Copies will be performed GPU-side directly into the output buffer +#ifdef USE_GPU_COPY + GLuint targetVBO = cloths[flagIndex].getVBO(); + btOpenGLInteropVertexBufferDescriptor *vertexBufferDescriptor = new btOpenGLInteropVertexBufferDescriptor(g_cqCommandQue, g_cxMainContext, targetVBO, 0, 8, 3, 8); +#else btCPUVertexBufferDescriptor *vertexBufferDescriptor = new btCPUVertexBufferDescriptor(reinterpret_cast< float* >(cloths[flagIndex].cpu_buffer), 0, 8, 3, 8); +#endif cloths[flagIndex].m_vertexBufferDescriptor = vertexBufferDescriptor; } @@ -402,40 +497,49 @@ void initBullet(void) -#ifndef BT_NO_PROFILE + btClock m_clock; -#endif //BT_NO_PROFILE void doFlags() { //float ms = getDeltaTimeMicroseconds(); -#ifndef BT_NO_PROFILE btScalar dt = (btScalar)m_clock.getTimeMicroseconds(); m_clock.reset(); -#else - btScalar dt = 1000000.f/60.f; -#endif ///step the simulation if( m_dynamicsWorld ) { m_dynamicsWorld->stepSimulation(dt/1000000.); + static int frameCount = 0; frameCount++; if (frameCount==100) { m_dynamicsWorld->stepSimulation(1./60.,0); -#ifndef BT_NO_PROFILE + + btDefaultSerializer* serializer = new btDefaultSerializer(); + m_dynamicsWorld->serialize(serializer); + + FILE* file = fopen("testFile.bullet","wb"); + fwrite(serializer->getBufferPointer(),serializer->getCurrentBufferSize(),1, file); + fclose(file); + CProfileManager::dumpAll(); -#endif //BT_NO_PROFILE } updatePhysicsWorld(); + + //m_dynamicsWorld->setDebugDrawer(&debugDraw); + //debugDraw.setDebugMode(btIDebugDraw::DBG_DrawWireframe); + //g_solver->copyBackToSoftBodies(); + + //m_dynamicsWorld->debugDrawWorld(); + } for( int flagIndex = 0; flagIndex < m_flags.size(); ++flagIndex ) { - g_solver->copySoftBodyToVertexBuffer( m_flags[flagIndex], cloths[flagIndex].m_vertexBufferDescriptor ); + g_softBodyOutput->copySoftBodyToVertexBuffer( m_flags[flagIndex], cloths[flagIndex].m_vertexBufferDescriptor ); cloths[flagIndex].draw(); } } @@ -444,9 +548,25 @@ void doFlags() int main(int argc, char *argv[]) { + + preInitGL(argc, argv); + glewInit(); + +#ifdef USE_GPU_COPY +#ifdef _WIN32 + HGLRC glCtx = wglGetCurrentContext(); +#else //!_WIN32 + GLXContext glCtx = glXGetCurrentContext(); +#endif //!_WIN32 + HDC glDC = wglGetCurrentDC(); + + initCL(glCtx, glDC); +#else initCL(); +#endif + cloths.resize(numFlags); for( int flagIndex = 0; flagIndex < numFlags; ++flagIndex ) @@ -457,8 +577,6 @@ int main(int argc, char *argv[]) initBullet(); m_dynamicsWorld->stepSimulation(1./60.,0); - preInitGL(argc, argv); - std::string flagTexs[] = { "atiFlag.bmp", "amdFlag.bmp", @@ -474,6 +592,16 @@ int main(int argc, char *argv[]) } goGL(); + + if( g_cpuSolver ) + delete g_cpuSolver; + if( g_openCLSolver ) + delete g_openCLSolver; + if( g_openCLSIMDSolver ) + delete g_openCLSIMDSolver; + if( g_softBodyOutput ) + delete g_softBodyOutput; + return 0; } diff --git a/Demos/OpenCLClothDemo/cloth.h b/Demos/OpenCLClothDemo/cloth.h index 386a8928d..dd2ddd7a9 100644 --- a/Demos/OpenCLClothDemo/cloth.h +++ b/Demos/OpenCLClothDemo/cloth.h @@ -49,6 +49,7 @@ class piece_of_cloth created = false; cpu_buffer = NULL; m_vertexBufferDescriptor = NULL; + clothVBO = 0; } bool created; @@ -63,6 +64,13 @@ class piece_of_cloth int height; GLuint m_texture; + + GLuint clothVBO; + + GLuint getVBO() + { + return clothVBO; + } void draw(void) { @@ -73,22 +81,41 @@ class piece_of_cloth glColor3f(0.0f, 1.0f, 1.0f); - glEnableClientState(GL_VERTEX_ARRAY); - //glEnableClientState(GL_NORMAL_ARRAY); - glEnableClientState(GL_TEXTURE_COORD_ARRAY); + int error = 0; + glBindBuffer(GL_ARRAY_BUFFER, clothVBO); +#ifndef USE_GPU_COPY + // Upload data to VBO + // Needed while we're not doing interop + glBufferData(GL_ARRAY_BUFFER, sizeof(vertex_struct)*width*height, &(cpu_buffer[0]), GL_DYNAMIC_DRAW); +#endif + + + glEnableClientState(GL_VERTEX_ARRAY); + glEnableClientState(GL_NORMAL_ARRAY); + glEnableClientState(GL_TEXTURE_COORD_ARRAY); + glBindTexture(GL_TEXTURE_2D, m_texture); - glVertexPointer( 3, GL_FLOAT, sizeof(vertex_struct), reinterpret_cast< GLvoid* >(&(cpu_buffer[0].pos[0])) ); - //glNormalPointer( 3, sizeof(vertex_struct), reinterpret_cast< GLvoid* >(&(cpu_buffer[0].normal[0])) ); - glTexCoordPointer( 2, GL_FLOAT, sizeof(vertex_struct), reinterpret_cast< GLvoid* >(&(cpu_buffer[0].texcoord[0])) ); + error = glGetError(); + + // VBO version + glVertexPointer( 3, GL_FLOAT, sizeof(vertex_struct), (const GLvoid *)0 ); + error = glGetError(); + glNormalPointer( GL_FLOAT, sizeof(vertex_struct), (const GLvoid *)(sizeof(float)*3) ); + error = glGetError(); + glTexCoordPointer( 2, GL_FLOAT, sizeof(vertex_struct), (const GLvoid *)(sizeof(float)*6) ); + error = glGetError(); + glDrawElements(GL_TRIANGLES, (height-1 )*(width-1)*3*2, GL_UNSIGNED_INT, indices); -// glDisableClientState(GL_NORMAL_ARRAY); + glDisableClientState(GL_NORMAL_ARRAY); glDisableClientState(GL_VERTEX_ARRAY); glDisableClientState(GL_TEXTURE_COORD_ARRAY); - + error = glGetError(); glBindTexture(GL_TEXTURE_2D, 0); + glBindBuffer(GL_ARRAY_BUFFER, 0); + error = glGetError(); } void create_texture(std::string filename) @@ -221,5 +248,16 @@ class piece_of_cloth indices[baseIndex+5] = x+width + y*width; } } + + // Construct VBO + glGenBuffers(1, &clothVBO); + glBindBuffer(GL_ARRAY_BUFFER, clothVBO); + // Do initial upload to ensure that the buffer exists on the device + glBufferData(GL_ARRAY_BUFFER, sizeof(vertex_struct)*width*height, &(cpu_buffer[0]), GL_DYNAMIC_DRAW); + int error = glGetError(); + glBindBuffer(GL_ARRAY_BUFFER, 0); + + + } }; diff --git a/Demos/OpenCLClothDemo/clstuff.cpp b/Demos/OpenCLClothDemo/clstuff.cpp index f6e6d4017..856d6b801 100644 --- a/Demos/OpenCLClothDemo/clstuff.cpp +++ b/Demos/OpenCLClothDemo/clstuff.cpp @@ -13,7 +13,7 @@ subject to the following restrictions: 3. This notice may not be removed or altered from any source distribution. */ -#include + #include "clstuff.h" #include "gl_win.h" @@ -27,9 +27,20 @@ cl_context g_cxMainContext; cl_device_id g_cdDevice; cl_command_queue g_cqCommandQue; -void initCL(void) +void initCL( void* glCtx, void* glDC ) { int ciErrNum = 0; + +#if defined(CL_PLATFORM_MINI_CL) + cl_device_type deviceType = CL_DEVICE_TYPE_CPU; +#elif defined(CL_PLATFORM_AMD) + cl_device_type deviceType = CL_DEVICE_TYPE_GPU; +#elif defined(CL_PLATFORM_NVIDIA) + cl_device_type deviceType = CL_DEVICE_TYPE_GPU; +#else + cl_device_type deviceType = CL_DEVICE_TYPE_CPU; +#endif + //g_cxMainContext = btOclCommon::createContextFromType(CL_DEVICE_TYPE_ALL, &ciErrNum); //g_cxMainContext = btOclCommon::createContextFromType(CL_DEVICE_TYPE_GPU, &ciErrNum); //g_cxMainContext = btOclCommon::createContextFromType(CL_DEVICE_TYPE_CPU, &ciErrNum); @@ -37,9 +48,9 @@ void initCL(void) //#ifdef USE_MINICL // g_cxMainContext = btOclCommon::createContextFromType(CL_DEVICE_TYPE_DEBUG, &ciErrNum); //#else - g_cxMainContext = btOclCommon::createContextFromType(CL_DEVICE_TYPE_ALL, &ciErrNum); + g_cxMainContext = btOclCommon::createContextFromType(deviceType, &ciErrNum, glCtx, glDC); //#endif - + oclCHECKERROR(ciErrNum, CL_SUCCESS); diff --git a/Demos/OpenCLClothDemo/clstuff.h b/Demos/OpenCLClothDemo/clstuff.h index 09f6313eb..83b0bbc22 100644 --- a/Demos/OpenCLClothDemo/clstuff.h +++ b/Demos/OpenCLClothDemo/clstuff.h @@ -4,7 +4,8 @@ - -void initCL(void); +// OpenCL initialization. +// Takes an optional GL context which, if passed, will create an interop-enabled CL context. +void initCL( void* glContext = 0, void* glDC = 0 ); #endif //__CLSTUFF_HDR__ \ No newline at end of file diff --git a/Demos/ParticlesOpenCL/shaders.cpp b/Demos/ParticlesOpenCL/shaders.cpp index 59f028c97..79b01f5c3 100644 --- a/Demos/ParticlesOpenCL/shaders.cpp +++ b/Demos/ParticlesOpenCL/shaders.cpp @@ -1,6 +1,6 @@ #define STRINGIFY(A) #A - -#ifdef USE_AMD_OPENCL +//use most up-to-date AMD Radeon drivers to make point sprites work +//see also http://forums.amd.com/devforum/messageview.cfm?catid=392&threadid=129431 // vertex shader const char *vertexShader = STRINGIFY( uniform float pointRadius; // point size in world space @@ -14,35 +14,7 @@ void main() posEye = vec3(gl_ModelViewMatrix * vec4(gl_Vertex.xyz, 1.0)); float dist = length(posEye); gl_PointSize = pointRadius * (pointScale / dist); -// gl_PointSize = 4.0; -//hack around latest AMD graphics cards having troubles with point sprite rendering -//the problem is still unresolved on the 5870 card, and results in a black screen -//see also http://forums.amd.com/devforum/messageview.cfm?catid=392&threadid=129431 - gl_TexCoord = gl_MultiTexCoord0; - gl_Position = gl_ModelViewProjectionMatrix * vec4(gl_Vertex.xyz, 1.0); - - gl_FrontColor = gl_Color; -} -); -#else -// vertex shader -const char *vertexShader = STRINGIFY( -uniform float pointRadius; // point size in world space -uniform float pointScale; // scale to calculate size in pixels -uniform float densityScale; -uniform float densityOffset; -varying vec3 posEye; -void main() -{ - // calculate window-space point size - posEye = vec3(gl_ModelViewMatrix * vec4(gl_Vertex.xyz, 1.0)); - float dist = length(posEye); - gl_PointSize = pointRadius * (pointScale / dist); -// gl_PointSize = 4.0; -//hack around latest AMD graphics cards having troubles with point sprite rendering -//the problem is still unresolved on the 5870 card, and results in a black screen -//see also http://forums.amd.com/devforum/messageview.cfm?catid=392&threadid=129431 gl_TexCoord[0] = gl_MultiTexCoord0; gl_Position = gl_ModelViewProjectionMatrix * vec4(gl_Vertex.xyz, 1.0); @@ -51,7 +23,7 @@ void main() } ); -#endif + // pixel shader for rendering points as shaded spheres const char *spherePixelShader = STRINGIFY( diff --git a/Demos/SharedOpenCL/btOclCommon.cpp b/Demos/SharedOpenCL/btOclCommon.cpp index 48fe105d7..c7bbe8df3 100644 --- a/Demos/SharedOpenCL/btOclCommon.cpp +++ b/Demos/SharedOpenCL/btOclCommon.cpp @@ -29,10 +29,11 @@ static char* spPlatformVendor = "Unknown Vendor"; #endif +#ifndef CL_PLATFORM_MINI_CL +#include "CL/cl_gl.h" +#endif - - -cl_context btOclCommon::createContextFromType(cl_device_type deviceType, cl_int* pErrNum) +cl_context btOclCommon::createContextFromType(cl_device_type deviceType, cl_int* pErrNum, void* pGLContext, void* pGLDC ) { cl_uint numPlatforms; cl_platform_id platform = NULL; @@ -76,12 +77,27 @@ cl_context btOclCommon::createContextFromType(cl_device_type deviceType, cl_int* * If we could find our platform, use it. Otherwise pass a NULL and get whatever the * implementation thinks we should be using. */ - cl_context_properties cps[3] = + cl_context_properties cps[7] = { CL_CONTEXT_PLATFORM, - (cl_context_properties)platform, - 0 + (cl_context_properties)platform, + 0, + 0, + 0, + 0, + 0 }; +#ifndef CL_PLATFORM_MINI_CL + // If we have a gl context then enable interop + if( pGLContext ) + { + cps[2] = CL_GL_CONTEXT_KHR; + cps[3] = (cl_context_properties)pGLContext; + cps[4] = CL_WGL_HDC_KHR; + cps[5] = (cl_context_properties)pGLDC; + } +#endif //CL_PLATFORM_MINI_CL + /* Use NULL for backward compatibility */ cl_context_properties* cprops = (NULL == platform) ? NULL : cps; cl_context retContext = clCreateContextFromType(cprops, diff --git a/Demos/SharedOpenCL/btOclCommon.h b/Demos/SharedOpenCL/btOclCommon.h index a4baa1dcb..ae3cd8344 100644 --- a/Demos/SharedOpenCL/btOclCommon.h +++ b/Demos/SharedOpenCL/btOclCommon.h @@ -34,7 +34,10 @@ subject to the following restrictions: class btOclCommon { public: - static cl_context createContextFromType(cl_device_type deviceType, cl_int* pErrNum); + // CL Context optionally takes a GL context. This is a generic type because we don't really want this code + // to have to understand GL types. + // It is a HGLRC in _WIN32 or a GLXContext otherwise. + static cl_context createContextFromType(cl_device_type deviceType, cl_int* pErrNum, void* pGLCtx = 0, void* pGLDC = 0); }; diff --git a/Demos/SharedOpenCL/btOclUtils.cpp b/Demos/SharedOpenCL/btOclUtils.cpp index 1ce0b9526..a8aca18a4 100644 --- a/Demos/SharedOpenCL/btOclUtils.cpp +++ b/Demos/SharedOpenCL/btOclUtils.cpp @@ -87,7 +87,32 @@ cl_device_id btOclGetMaxFlopsDev(cl_context cxMainContext) cl_uint clock_frequency; clGetDeviceInfo(cdDevices[current_device], CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(clock_frequency), &clock_frequency, NULL); - max_flops = compute_units * clock_frequency; + cl_device_type device_type; + clGetDeviceInfo(cdDevices[current_device], CL_DEVICE_TYPE, sizeof(device_type), &device_type, NULL); + + int SIMDmultiplier = 1; + + if( device_type == CL_DEVICE_TYPE_CPU ) + { + // For simplicity assume that the CPU is running single SSE instructions + // This will of course depend on the kernel + SIMDmultiplier = 4; + } else if( device_type == CL_DEVICE_TYPE_GPU ) { + // Approximation to GPU compute power + // As long as this beats the CPU number that's the important thing, really +#if defined(CL_PLATFORM_AMD) + // 16 processing elements, 5 ALUs each + SIMDmultiplier = 80; +#elif defined(CL_PLATFORM_NVIDIA) + // 8 processing elements, dual issue - pre-Fermi at least + SIMDmultiplier = 16; +#else + SIMDmultiplier = 1; +#endif + } + + + max_flops = compute_units * clock_frequency * SIMDmultiplier; ++current_device; while( current_device < device_count ) diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/CPU/btSoftBodySolverData.h b/src/BulletMultiThreaded/GpuSoftBodySolvers/CPU/btSoftBodySolverData.h index df8d9e226..eb4e98b08 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/CPU/btSoftBodySolverData.h +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/CPU/btSoftBodySolverData.h @@ -435,11 +435,16 @@ public: m_vertexTriangleCount[vertexIndex] = 0; } - /** Create numVertices new vertices for cloth clothIdentifier */ - void createVertices( int numVertices, int clothIdentifier ) + /** + * Create numVertices new vertices for cloth clothIdentifier + * maxVertices allows a buffer zone of extra vertices for alignment or tearing reasons. + */ + void createVertices( int numVertices, int clothIdentifier, int maxVertices = 0 ) { int previousSize = m_vertexPosition.size(); - int newSize = previousSize + numVertices; + if( maxVertices == 0 ) + maxVertices = numVertices; + int newSize = previousSize + maxVertices; // Resize all the arrays that store vertex data m_clothIdentifier.resize( newSize ); @@ -454,6 +459,8 @@ public: for( int vertexIndex = previousSize; vertexIndex < newSize; ++vertexIndex ) m_clothIdentifier[vertexIndex] = clothIdentifier; + for( int vertexIndex = (previousSize + numVertices); vertexIndex < newSize; ++vertexIndex ) + m_clothIdentifier[vertexIndex] = -1; } // Get and set methods in header so they can be inlined @@ -466,6 +473,11 @@ public: return m_vertexPosition[vertexIndex]; } + Vectormath::Aos::Point3 getPosition( int vertexIndex ) const + { + return m_vertexPosition[vertexIndex]; + } + /** * Return a reference to the previous position of vertex vertexIndex as stored on the host. */ @@ -498,6 +510,11 @@ public: return m_vertexNormal[vertexIndex]; } + Vectormath::Aos::Vector3 getNormal( int vertexIndex ) const + { + return m_vertexNormal[vertexIndex]; + } + /** * Return a reference to the inverse mass of vertex vertexIndex as stored on the host. */ diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/CPU/btSoftBodySolver_CPU.cpp b/src/BulletMultiThreaded/GpuSoftBodySolvers/CPU/btSoftBodySolver_CPU.cpp index a7ff177f6..51a24baff 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/CPU/btSoftBodySolver_CPU.cpp +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/CPU/btSoftBodySolver_CPU.cpp @@ -74,12 +74,46 @@ static Vectormath::Aos::Transform3 toTransform3( const btTransform &transform ) return outTransform; } - - - -void btCPUSoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * > &softBodies ) +void btCPUSoftBodySolver::btAcceleratedSoftBodyInterface::updateBounds( const btVector3 &lowerBound, const btVector3 &upperBound ) { - if( m_softBodySet.size() != softBodies.size() ) + float scalarMargin = this->getSoftBody()->getCollisionShape()->getMargin(); + btVector3 vectorMargin( scalarMargin, scalarMargin, scalarMargin ); + m_softBody->m_bounds[0] = lowerBound - vectorMargin; + m_softBody->m_bounds[1] = upperBound + vectorMargin; +} + + +void btCPUSoftBodySolver::copyBackToSoftBodies() +{ + // Loop over soft bodies, copying all the vertex positions back for each body in turn + for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) + { + btAcceleratedSoftBodyInterface *softBodyInterface = m_softBodySet[ softBodyIndex ]; + btSoftBody *softBody = softBodyInterface->getSoftBody(); + + int firstVertex = softBodyInterface->getFirstVertex(); + int numVertices = softBodyInterface->getNumVertices(); + + // Copy vertices from solver back into the softbody + for( int vertex = 0; vertex < numVertices; ++vertex ) + { + using Vectormath::Aos::Point3; + Point3 vertexPosition( getVertexData().getVertexPositions()[firstVertex + vertex] ); + + softBody->m_nodes[vertex].m_x.setX( vertexPosition.getX() ); + softBody->m_nodes[vertex].m_x.setY( vertexPosition.getY() ); + softBody->m_nodes[vertex].m_x.setZ( vertexPosition.getZ() ); + + softBody->m_nodes[vertex].m_n.setX( vertexPosition.getX() ); + softBody->m_nodes[vertex].m_n.setY( vertexPosition.getY() ); + softBody->m_nodes[vertex].m_n.setZ( vertexPosition.getZ() ); + } + } +} // btCPUSoftBodySolver::copyBackToSoftBodies + +void btCPUSoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * > &softBodies , bool forceUpdate ) +{ + if( forceUpdate || m_softBodySet.size() != softBodies.size() ) { // Have a change in the soft body set so update, reloading all the data getVertexData().clear(); @@ -104,6 +138,7 @@ void btCPUSoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * > &softBo m_perClothLiftFactor.push_back( softBody->m_cfg.kLF ); m_perClothDragFactor.push_back( softBody->m_cfg.kDG ); m_perClothMediumDensity.push_back(softBody->getWorldInfo()->air_density); + m_perClothCollisionObjects.push_back( CollisionObjectIndices(-1, -1) ); // Add space for new vertices and triangles in the default solver for now // TODO: Include space here for tearing too later @@ -385,51 +420,72 @@ void btCPUSoftBodySolver::updateConstants( float timeStep ) } } // btCPUSoftBodySolver::updateConstants + + + +void btCPUSoftBodySolver::updateBounds() +{ + using Vectormath::Aos::Point3; + + for( int clothIndex = 0; clothIndex < m_softBodySet.size(); ++clothIndex ) + { + btAcceleratedSoftBodyInterface *currentCloth = m_softBodySet[clothIndex]; + btVector3 startBound(FLT_MAX, FLT_MAX, FLT_MAX); + btVector3 endBound(FLT_MIN, FLT_MIN, FLT_MIN); + + const int startVertex = currentCloth->getFirstVertex(); + const int numVertices = currentCloth->getNumVertices(); + + int endVertex = startVertex + numVertices; + for(int vertexIndex = startVertex; vertexIndex < endVertex; ++vertexIndex) + { + btVector3 vertexPosition( m_vertexData.getVertexPositions()[vertexIndex].getX(), m_vertexData.getVertexPositions()[vertexIndex].getY(), m_vertexData.getVertexPositions()[vertexIndex].getZ() ); + startBound.setX( btMin( startBound.getX(), vertexPosition.getX() ) ); + startBound.setY( btMin( startBound.getY(), vertexPosition.getY() ) ); + startBound.setZ( btMin( startBound.getZ(), vertexPosition.getZ() ) ); + + endBound.setX( btMax( endBound.getX(), vertexPosition.getX() ) ); + endBound.setY( btMax( endBound.getY(), vertexPosition.getY() ) ); + endBound.setZ( btMax( endBound.getZ(), vertexPosition.getZ() ) ); + } + + m_softBodySet[clothIndex]->updateBounds( startBound, endBound ); + } +} + + +class btCPUSB_QuickSortCompare +{ + public: + + bool operator() ( const btCPUCollisionShapeDescription& a, const btCPUCollisionShapeDescription& b ) + { + return ( a.softBodyIdentifier < b.softBodyIdentifier ); + } +}; + /** * Sort the collision object details array and generate indexing into it for the per-cloth collision object array. */ void btCPUSoftBodySolver::prepareCollisionConstraints() { - // First do a simple radix sort on the collision objects + // First do a simple sort on the collision objects btAlignedObjectArray numObjectsPerClothPrefixSum; btAlignedObjectArray numObjectsPerCloth; numObjectsPerCloth.resize( m_softBodySet.size(), 0 ); numObjectsPerClothPrefixSum.resize( m_softBodySet.size(), 0 ); - btAlignedObjectArray< CollisionShapeDescription > m_collisionObjectDetailsCopy(m_collisionObjectDetails); - // Count and prefix sum number of previous cloths - for( int collisionObject = 0; collisionObject < m_collisionObjectDetailsCopy.size(); ++collisionObject ) - { - CollisionShapeDescription &shapeDescription( m_collisionObjectDetailsCopy[collisionObject] ); - ++numObjectsPerClothPrefixSum[shapeDescription.softBodyIdentifier]; - } - int sum = 0; - for( int cloth = 0; cloth < m_softBodySet.size(); ++cloth ) - { - int currentValue = numObjectsPerClothPrefixSum[cloth]; - numObjectsPerClothPrefixSum[cloth] = sum; - sum += currentValue; - } - // Move into the target array - for( int collisionObject = 0; collisionObject < m_collisionObjectDetailsCopy.size(); ++collisionObject ) - { - CollisionShapeDescription &shapeDescription( m_collisionObjectDetailsCopy[collisionObject] ); - int clothID = shapeDescription.softBodyIdentifier; - int newLocation = numObjectsPerClothPrefixSum[clothID] + numObjectsPerCloth[clothID]; - numObjectsPerCloth[shapeDescription.softBodyIdentifier]++; - m_collisionObjectDetails[newLocation] = shapeDescription; - } - for( int collisionObject = 0; collisionObject < m_collisionObjectDetailsCopy.size(); ++collisionObject ) - { - CollisionShapeDescription &shapeDescription( m_collisionObjectDetails[collisionObject] ); - } + if (!m_perClothCollisionObjects.size()) + return; + + m_collisionObjectDetails.quickSort( btCPUSB_QuickSortCompare() ); // Generating indexing for perClothCollisionObjects - // First clear the previous values + // First clear the previous values with the "no collision object for cloth" constant for( int clothIndex = 0; clothIndex < m_perClothCollisionObjects.size(); ++clothIndex ) { - m_perClothCollisionObjects[clothIndex].firstObject = 0; - m_perClothCollisionObjects[clothIndex].endObject = 0; + m_perClothCollisionObjects[clothIndex].firstObject = -1; + m_perClothCollisionObjects[clothIndex].endObject = -1; } int currentCloth = 0; int startIndex = 0; @@ -445,8 +501,12 @@ void btCPUSoftBodySolver::prepareCollisionConstraints() currentCloth = nextCloth; startIndex = collisionObject; } - } - //m_perClothCollisionObjects + } + + // And update last cloth + m_perClothCollisionObjects[currentCloth].firstObject = startIndex; + m_perClothCollisionObjects[currentCloth].endObject = m_collisionObjectDetails.size(); + } // prepareCollisionConstraints @@ -478,83 +538,9 @@ void btCPUSoftBodySolver::solveConstraints( float solverdt ) } -#if 0 + prepareCollisionConstraints(); - // Solve collision constraints - // Very simple solver that pushes the vertex out of collision imposters for now - // to test integration with the broad phase code. - // May also want to put this into position solver loop every n iterations depending on - // how it behaves - for( int clothIndex = 0; clothIndex < m_softBodySet.size(); ++clothIndex ) - { - btAcceleratedSoftBodyInterface *currentCloth = m_softBodySet[clothIndex]; - - const int startVertex = currentCloth->getFirstVertex(); - const int numVertices = currentCloth->getNumVertices(); - int endVertex = startVertex + numVertices; - - int startObject = m_perClothCollisionObjects[clothIndex].firstObject; - int endObject = m_perClothCollisionObjects[clothIndex].endObject; - - for( int collisionObject = startObject; collisionObject < endObject; ++collisionObject ) - { - CollisionShapeDescription &shapeDescription( m_collisionObjectDetails[collisionObject] ); - - if( shapeDescription.collisionShapeType == CAPSULE_SHAPE_PROXYTYPE ) - { - using namespace Vectormath::Aos; - - float capsuleHalfHeight = shapeDescription.shapeInformation.capsule.halfHeight; - float capsuleRadius = shapeDescription.shapeInformation.capsule.radius; - Transform3 worldTransform = shapeDescription.shapeTransform; - for( int vertexIndex = startVertex; vertexIndex < endVertex; ++vertexIndex ) - { - Point3 vertex( m_vertexData.getPosition( vertexIndex ) ); - Point3 c1(0.f, -capsuleHalfHeight, 0.f); - Point3 c2(0.f, +capsuleHalfHeight, 0.f); - Point3 worldC1 = worldTransform * c1; - Point3 worldC2 = worldTransform * c2; - Vector3 segment = worldC2 - worldC1; - - // compute distance of tangent to vertex along line segment in capsule - float distanceAlongSegment = -( dot( worldC1 - vertex, segment ) / lengthSqr(segment) ); - - Point3 closestPoint = (worldC1 + segment * distanceAlongSegment); - float distanceFromLine = length(vertex - closestPoint); - float distanceFromC1 = length(worldC1 - vertex); - float distanceFromC2 = length(worldC2 - vertex); - - // Final distance from collision, point to push from, direction to push in - // for impulse force - float distance; - Point3 sourcePoint; - Vector3 pushVector; - if( distanceAlongSegment < 0 ) - { - distance = distanceFromC1; - sourcePoint = worldC1; - pushVector = normalize(vertex - worldC1); - } else if( distanceAlongSegment > 1.f ) { - distance = distanceFromC1; - sourcePoint = worldC1; - pushVector = normalize(vertex - worldC1); - } else { - distance = distanceFromLine; - sourcePoint = closestPoint; - pushVector = normalize(vertex - closestPoint); - } - - // For now just update vertex position by moving to radius distance along the push vector - // Could use this as the basis for simple vector distance constraint for the point later, possibly? - // That way in the main solver loop all shape types could be the same... though when - // we need to apply bi-directionally it becomes more complicated - m_vertexData.getPosition( vertexIndex ) = closestPoint + capsuleRadius * pushVector; - } - } - } - } -#endif for( int iteration = 0; iteration < m_numberOfVelocityIterations ; ++iteration ) { @@ -622,6 +608,8 @@ void btCPUSoftBodySolver::solveConstraints( float solverdt ) } } } + + // Clear forces so that friction is applied correctly for( int clothIndex = 0; clothIndex < m_softBodySet.size(); ++clothIndex ) { btAcceleratedSoftBodyInterface *currentCloth = m_softBodySet[clothIndex]; @@ -637,23 +625,201 @@ void btCPUSoftBodySolver::solveConstraints( float solverdt ) float velocityCorrectionCoefficient = m_perClothVelocityCorrectionCoefficient[clothIndex]; float isolverDt = 1.f/solverdt; - if( m_numberOfVelocityIterations > 0 ) + for(int vertexIndex = startVertex; vertexIndex < lastVertex; ++vertexIndex) { - for(int vertexIndex = startVertex; vertexIndex < lastVertex; ++vertexIndex) - { - m_vertexData.getVelocity(vertexIndex) += (m_vertexData.getPosition(vertexIndex) - m_vertexData.getPreviousPosition(vertexIndex)) * velocityCorrectionCoefficient * isolverDt; - m_vertexData.getVelocity(vertexIndex) *= velocityCoefficient; - m_vertexData.getForceAccumulator( vertexIndex ) = Vector3(0.f, 0.f, 0.f); - } - } else { - // If we didn't compute the velocity iteratively then we compute it purely based on the position change - for(int vertexIndex = startVertex; vertexIndex < lastVertex; ++vertexIndex) - { - m_vertexData.getVelocity(vertexIndex) = (m_vertexData.getPosition(vertexIndex) - m_vertexData.getPreviousPosition(vertexIndex)) * velocityCoefficient * isolverDt; - m_vertexData.getForceAccumulator( vertexIndex ) = Vector3(0.f, 0.f, 0.f); - } + m_vertexData.getForceAccumulator( vertexIndex ) = Vector3(0.f, 0.f, 0.f); } } + + + + + // Solve collision constraints + // Very simple solver that pushes the vertex out of collision imposters for now + // to test integration with the broad phase code. + // May also want to put this into position solver loop every n iterations depending on + // how it behaves + for( int clothIndex = 0; clothIndex < m_softBodySet.size(); ++clothIndex ) + { + btAcceleratedSoftBodyInterface *currentCloth = m_softBodySet[clothIndex]; + + float clothFriction = currentCloth->getSoftBody()->getFriction(); + + const int startVertex = currentCloth->getFirstVertex(); + const int numVertices = currentCloth->getNumVertices(); + int endVertex = startVertex + numVertices; + + float velocityCoefficient = (1.f - m_perClothDampingFactor[clothIndex]); + float velocityCorrectionCoefficient = m_perClothVelocityCorrectionCoefficient[clothIndex]; + float isolverDt = 1.f/solverdt; + + int startObject = m_perClothCollisionObjects[clothIndex].firstObject; + int endObject = m_perClothCollisionObjects[clothIndex].endObject; + + if( endObject == startObject ) + { + // No collisions so just do the force update + for(int vertexIndex = startVertex; vertexIndex < endVertex; ++vertexIndex) + { + m_vertexData.getForceAccumulator( vertexIndex ) = Vector3(0.f, 0.f, 0.f); + } + + // Recompute velocity based on updated position inclusive of drift + for(int vertexIndex = startVertex; vertexIndex < endVertex; ++vertexIndex) + { + m_vertexData.getVelocity(vertexIndex) = (m_vertexData.getPosition(vertexIndex) - m_vertexData.getPreviousPosition(vertexIndex)) * velocityCoefficient * isolverDt; + } + } else { + + for( int collisionObject = startObject; collisionObject < endObject; ++collisionObject ) + { + btCPUCollisionShapeDescription &shapeDescription( m_collisionObjectDetails[collisionObject] ); + + float colliderFriction = shapeDescription.friction; + + if( shapeDescription.collisionShapeType == CAPSULE_SHAPE_PROXYTYPE ) + { + using namespace Vectormath::Aos; + + float capsuleHalfHeight = shapeDescription.shapeInformation.capsule.halfHeight; + float capsuleRadius = shapeDescription.shapeInformation.capsule.radius; + int capsuleUpAxis = shapeDescription.shapeInformation.capsule.upAxis; + float capsuleMargin = shapeDescription.margin; + + Transform3 worldTransform = shapeDescription.shapeTransform; + + // As this is a GPU comparison solver just iterate over the vertices + for( int vertexIndex = startVertex; vertexIndex < endVertex; ++vertexIndex ) + { + // Clear force for vertex first + m_vertexData.getForceAccumulator( vertexIndex ) = Vector3(0.f, 0.f, 0.f); + + Point3 vertex( m_vertexData.getPosition( vertexIndex ) ); + + // Correctly define the centerline depending on the upAxis + Point3 c1(0.f, 0.f, 0.f); + Point3 c2(0.f, 0.f, 0.f); + if( capsuleUpAxis == 0 ) { + c1.setX(-capsuleHalfHeight); + c2.setX(capsuleHalfHeight); + } else if( capsuleUpAxis == 1 ) { + c1.setY(-capsuleHalfHeight); + c2.setY(capsuleHalfHeight); + } else { + c1.setZ(-capsuleHalfHeight); + c2.setZ(capsuleHalfHeight); + } + + Point3 worldC1 = worldTransform * c1; + Point3 worldC2 = worldTransform * c2; + Vector3 segment = worldC2 - worldC1; + + // compute distance of tangent to vertex along line segment in capsule + float distanceAlongSegment = -( dot( worldC1 - vertex, segment ) / lengthSqr(segment) ); + + Point3 closestPoint = (worldC1 + segment * distanceAlongSegment); + float distanceFromLine = length(vertex - closestPoint); + float distanceFromC1 = length(worldC1 - vertex); + float distanceFromC2 = length(worldC2 - vertex); + + // Final distance from collision, point to push from, direction to push in + // for impulse force + float distance; + Point3 sourcePoint; + Vector3 normalVector; + if( distanceAlongSegment < 0 ) + { + distance = distanceFromC1; + sourcePoint = worldC1; + normalVector = normalize(vertex - worldC1); + } else if( distanceAlongSegment > 1.f ) { + distance = distanceFromC2; + sourcePoint = worldC2; + normalVector = normalize(vertex - worldC2); + } else { + distance = distanceFromLine; + sourcePoint = closestPoint; + normalVector = normalize(vertex - closestPoint); + } + + Vector3 colliderLinearVelocity( shapeDescription.linearVelocity ); + Vector3 colliderAngularVelocity( shapeDescription.angularVelocity ); + Vector3 velocityOfSurfacePoint = colliderLinearVelocity + cross(colliderAngularVelocity, Vector3(vertex) - worldTransform.getTranslation()); + + float minDistance = capsuleRadius + capsuleMargin; + bool collided = false; + + if( distance < minDistance ) + { + // Project back to surface along normal + Vectormath::Aos::Point3 sourcePos = m_vertexData.getPosition( vertexIndex ); + Vectormath::Aos::Vector3 posChange = (minDistance - distance)*normalVector*0.9; + //if( length(posChange) > 1 ) + // std::cerr << "Poschange: " << length(posChange) << "\n"; + + Vectormath::Aos::Point3 newPos = sourcePos + posChange; + m_vertexData.getPosition( vertexIndex ) = newPos; + //m_vertexData.getPosition( vertexIndex ) = m_vertexData.getPosition( vertexIndex ) + (minDistance - distance)*normalVector*0.9; + + // Experiment with moving back along source vector. + // Removes all ability to slide because it projects back along the vector length so it would undo lateral movement. + // TODO: This isn't quite right because we should take the movement of the collider into account as well + /*Vector3 incomingMoveVector( normalize(m_vertexData.getPreviousPosition(vertexIndex) - m_vertexData.getPosition(vertexIndex)) ); + Vector3 normalDirectionMoveOut( (minDistance - distance)*normalVector*0.9 ); + float distanceOnIncomingVector = dot(normalDirectionMoveOut, incomingMoveVector); + Vector3 vectorCorrection( distanceOnIncomingVector*incomingMoveVector ); + m_vertexData.getPosition( vertexIndex ) = m_vertexData.getPosition( vertexIndex ) + vectorCorrection;*/ + + + collided = true; + } + + // Update velocity of vertex based on position + m_vertexData.getVelocity(vertexIndex) = (m_vertexData.getPosition(vertexIndex) - m_vertexData.getPreviousPosition(vertexIndex)) * velocityCoefficient * isolverDt; + + // If we collided before we are on the surface so have friction + if( collided ) + { + // Compute friction + + // TODO: Just vertex velocity not enough, really we need the velocity + // relative to closest point on the surface of the collider + Vector3 vertexVelocity( m_vertexData.getVelocity(vertexIndex) ); + Vector3 relativeVelocity( vertexVelocity - velocityOfSurfacePoint ); + + + // First compute vectors for plane perpendicular to normal vector + // Cross any vector with normal vector first then cross the normal with it again + Vector3 p1(normalize(cross(normalVector, segment))); + Vector3 p2(normalize(cross(p1, normalVector))); + // Full friction is sum of velocities in each direction of plane. + Vector3 frictionVector(p1*dot(relativeVelocity, p1) + p2*dot(relativeVelocity, p2)); + + // Real friction is peak friction corrected by friction coefficients. + frictionVector = frictionVector*(colliderFriction*clothFriction); + + float approachSpeed = dot( relativeVelocity, normalVector ); + + // For now just update vertex position by moving to radius distance along the push vector + // Could use this as the basis for simple vector distance constraint for the point later, possibly? + // That way in the main solver loop all shape types could be the same... though when + // we need to apply bi-directionally it becomes more complicated + + // Add friction vector to the force accumulator + Vector3 ¤tForce( m_vertexData.getForceAccumulator( vertexIndex ) ); + + // Only apply if the vertex is moving towards the object to reduce jitter error + if( approachSpeed <= 0.0 ) + currentForce -= frictionVector; + } + } + } + } // for( int collisionObject = startObject; collisionObject < endObject; ++collisionObject ) + } // if( endObject == startObject ) + } + + + } // btCPUSoftBodySolver::solveConstraints @@ -669,14 +835,37 @@ btCPUSoftBodySolver::btAcceleratedSoftBodyInterface *btCPUSoftBodySolver::findSo return 0; } -void btCPUSoftBodySolver::copySoftBodyToVertexBuffer( const btSoftBody * const softBody, btVertexBufferDescriptor *vertexBuffer ) +const btCPUSoftBodySolver::btAcceleratedSoftBodyInterface * const btCPUSoftBodySolver::findSoftBodyInterface( const btSoftBody* const softBody ) const +{ + for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) + { + const btAcceleratedSoftBodyInterface *const softBodyInterface = m_softBodySet[softBodyIndex]; + if( softBodyInterface->getSoftBody() == softBody ) + return softBodyInterface; + } + return 0; +} + +int btCPUSoftBodySolver::findSoftBodyIndex( const btSoftBody* const softBody ) +{ + for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) + { + btAcceleratedSoftBodyInterface *softBodyInterface = m_softBodySet[softBodyIndex]; + if( softBodyInterface->getSoftBody() == softBody ) + return softBodyIndex; + } + return 1; +} + +void btSoftBodySolverOutputCPUtoCPU::copySoftBodyToVertexBuffer( const btSoftBody * const softBody, btVertexBufferDescriptor *vertexBuffer ) { // Currently only support CPU output buffers - // TODO: check for DX11 buffers. Take all offsets into the same DX11 buffer - // and use them together on a single kernel call if possible by setting up a - // per-cloth target buffer array for the copy kernel. - - btAcceleratedSoftBodyInterface *currentCloth = findSoftBodyInterface( softBody ); + + const btSoftBodySolver *solver = softBody->getSoftBodySolver(); + btAssert( solver->getSolverType() == btSoftBodySolver::CPU_SOLVER ); + const btCPUSoftBodySolver *cpuSolver = static_cast< const btCPUSoftBodySolver * >( solver ); + const btCPUSoftBodySolver::btAcceleratedSoftBodyInterface * const currentCloth = cpuSolver->findSoftBodyInterface( softBody ); + const btSoftBodyVertexData &vertexData( cpuSolver->m_vertexData ); if( vertexBuffer->getBufferType() == btVertexBufferDescriptor::CPU_BUFFER ) { @@ -693,7 +882,7 @@ void btCPUSoftBodySolver::copySoftBodyToVertexBuffer( const btSoftBody * const s for( int vertexIndex = firstVertex; vertexIndex < lastVertex; ++vertexIndex ) { - Vectormath::Aos::Point3 position = m_vertexData.getPosition(vertexIndex); + Vectormath::Aos::Point3 position = vertexData.getPosition(vertexIndex); *(vertexPointer + 0) = position.getX(); *(vertexPointer + 1) = position.getY(); *(vertexPointer + 2) = position.getZ(); @@ -708,40 +897,58 @@ void btCPUSoftBodySolver::copySoftBodyToVertexBuffer( const btSoftBody * const s for( int vertexIndex = firstVertex; vertexIndex < lastVertex; ++vertexIndex ) { - Vectormath::Aos::Vector3 normal = m_vertexData.getNormal(vertexIndex); + Vectormath::Aos::Vector3 normal = vertexData.getNormal(vertexIndex); *(normalPointer + 0) = normal.getX(); *(normalPointer + 1) = normal.getY(); *(normalPointer + 2) = normal.getZ(); normalPointer += normalStride; } } + } else { + btAssert( 0=="Invalid vertex buffer descriptor used in CPU output." ); } } // btCPUSoftBodySolver::outputToVertexBuffers - - -void btCPUSoftBodySolver::addCollisionObjectForSoftBody( int clothIndex, btCollisionObject *collisionObject ) +void btCPUSoftBodySolver::processCollision( btSoftBody*, btSoftBody *) { - btCollisionShape *collisionShape = collisionObject->getCollisionShape(); - int shapeType = collisionShape->getShapeType(); - if( shapeType == CAPSULE_SHAPE_PROXYTYPE ) - { - // Add to the list of expected collision objects - CollisionShapeDescription newCollisionShapeDescription; - newCollisionShapeDescription.softBodyIdentifier = clothIndex; - newCollisionShapeDescription.collisionShapeType = shapeType; - newCollisionShapeDescription.shapeTransform = toTransform3(collisionObject->getWorldTransform()); - btCapsuleShape *capsule = static_cast( collisionShape ); - newCollisionShapeDescription.shapeInformation.capsule.radius = capsule->getRadius(); - newCollisionShapeDescription.shapeInformation.capsule.halfHeight = capsule->getHalfHeight(); - m_collisionObjectDetails.push_back( newCollisionShapeDescription ); - // TODO: In the collision function, sort the above array on the clothIndex and generate the start and end indices - } else { - btAssert("Unsupported collision shape type\n"); - } } +// Add the collision object to the set to deal with for a particular soft body +void btCPUSoftBodySolver::processCollision( btSoftBody *softBody, btCollisionObject* collisionObject ) +{ + int softBodyIndex = findSoftBodyIndex( softBody ); + + if( softBodyIndex >= 0 ) + { + btCollisionShape *collisionShape = collisionObject->getCollisionShape(); + float friction = collisionObject->getFriction(); + int shapeType = collisionShape->getShapeType(); + if( shapeType == CAPSULE_SHAPE_PROXYTYPE ) + { + // Add to the list of expected collision objects + btCPUCollisionShapeDescription newCollisionShapeDescription; + newCollisionShapeDescription.softBodyIdentifier = softBodyIndex; + newCollisionShapeDescription.collisionShapeType = shapeType; + newCollisionShapeDescription.shapeTransform = toTransform3(collisionObject->getWorldTransform()); + btCapsuleShape *capsule = static_cast( collisionShape ); + newCollisionShapeDescription.shapeInformation.capsule.radius = capsule->getRadius(); + newCollisionShapeDescription.shapeInformation.capsule.halfHeight = capsule->getHalfHeight(); + newCollisionShapeDescription.shapeInformation.capsule.upAxis = capsule->getUpAxis(); + newCollisionShapeDescription.margin = capsule->getMargin(); + newCollisionShapeDescription.friction = friction; + btRigidBody* body = static_cast< btRigidBody* >( collisionObject ); + newCollisionShapeDescription.linearVelocity = toVector3(body->getLinearVelocity()); + newCollisionShapeDescription.angularVelocity = toVector3(body->getAngularVelocity()); + m_collisionObjectDetails.push_back( newCollisionShapeDescription ); + } else { + btAssert("Unsupported collision shape type\n"); + } + } else { + btAssert("Unknown soft body"); + } +} // btCPUSoftBodySolver::processCollision + void btCPUSoftBodySolver::predictMotion( float timeStep ) { @@ -760,6 +967,12 @@ void btCPUSoftBodySolver::predictMotion( float timeStep ) // Itegrate motion for all soft bodies dealt with by the solver integrate( timeStep * getTimeScale() ); + + // Update bounds + // Will update the bounds for all softBodies being dealt with by the solver and + // set the values in the btSoftBody object + updateBounds(); + // End prediction work for solvers } diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/CPU/btSoftBodySolver_CPU.h b/src/BulletMultiThreaded/GpuSoftBodySolvers/CPU/btSoftBodySolver_CPU.h index 18a749ea4..25d04cd9b 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/CPU/btSoftBodySolver_CPU.h +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/CPU/btSoftBodySolver_CPU.h @@ -21,7 +21,37 @@ subject to the following restrictions: #include "BulletSoftBody/btSoftBodySolverVertexBuffer.h" #include "BulletMultiThreaded/GpuSoftBodySolvers/CPU/btSoftBodySolverData.h" +struct btCPUCollisionShapeDescription +{ + int softBodyIdentifier; + int collisionShapeType; + Vectormath::Aos::Transform3 shapeTransform; + union + { + struct Sphere + { + float radius; + } sphere; + struct Capsule + { + float radius; + float halfHeight; + int upAxis; + } capsule; + } shapeInformation; + + float margin; + float friction; + Vectormath::Aos::Vector3 linearVelocity; + Vectormath::Aos::Vector3 angularVelocity; + btCPUCollisionShapeDescription() + { + collisionShapeType = 0; + margin = 0; + friction = 0; + } +}; class btCPUSoftBodySolver : public btSoftBodySolver { @@ -30,28 +60,22 @@ protected: * Entry in the collision shape array. * Specifies the shape type, the transform matrix and the necessary details of the collisionShape. */ - struct CollisionShapeDescription - { - int softBodyIdentifier; - int collisionShapeType; - Vectormath::Aos::Transform3 shapeTransform; - union - { - struct Sphere - { - float radius; - } sphere; - struct Capsule - { - float radius; - float halfHeight; - } capsule; - } shapeInformation; - CollisionShapeDescription() + + // Public because output classes need it. This is a better encapsulation to break in the short term + // Than having the solvers themselves need dependencies on DX, CL etc unnecessarily +public: + + struct CollisionObjectIndices + { + CollisionObjectIndices( int f, int e ) { - collisionShapeType = 0; + firstObject = f; + endObject = e; } + + int firstObject; + int endObject; }; /** @@ -99,36 +123,41 @@ protected: m_maxLinks = 0; m_numLinks = 0; } - int getNumVertices() + int getNumVertices() const { return m_numVertices; } - int getNumTriangles() + int getNumTriangles() const { return m_numTriangles; } - int getMaxVertices() + int getMaxVertices() const { return m_maxVertices; } - int getMaxTriangles() + int getMaxTriangles() const { return m_maxTriangles; } - int getFirstVertex() + int getFirstVertex() const { return m_firstVertex; } - int getFirstTriangle() + int getFirstTriangle() const { return m_firstTriangle; } + /** + * Update the bounds in the btSoftBody object + */ + void updateBounds( const btVector3 &lowerBound, const btVector3 &upperBound ); + // TODO: All of these set functions will have to do checks and // update the world because restructuring of the arrays will be necessary // Reasonable use of "friend"? @@ -177,17 +206,17 @@ protected: m_firstLink = firstLink; } - int getMaxLinks() + int getMaxLinks() const { return m_maxLinks; } - int getNumLinks() + int getNumLinks() const { return m_numLinks; } - int getFirstLink() + int getFirstLink() const { return m_firstLink; } @@ -197,47 +226,20 @@ protected: return m_softBody; } - #if 0 - void setAcceleration( Vectormath::Aos::Vector3 acceleration ) + const btSoftBody* const getSoftBody() const { - m_currentSolver->setPerClothAcceleration( m_clothIdentifier, acceleration ); + return m_softBody; } - - void setWindVelocity( Vectormath::Aos::Vector3 windVelocity ) - { - m_currentSolver->setPerClothWindVelocity( m_clothIdentifier, windVelocity ); - } - - /** - * Set the density of the air in which the cloth is situated. - */ - void setAirDensity( btScalar density ) - { - m_currentSolver->setPerClothMediumDensity( m_clothIdentifier, static_cast(density) ); - } - - /** - * Add a collision object to this soft body. - */ - void addCollisionObject( btCollisionObject *collisionObject ) - { - m_currentSolver->addCollisionObjectForSoftBody( m_clothIdentifier, collisionObject ); - } - #endif }; - - - struct CollisionObjectIndices - { - int firstObject; - int endObject; - }; - - - + btSoftBodyLinkData m_linkData; btSoftBodyVertexData m_vertexData; btSoftBodyTriangleData m_triangleData; + +protected: + + + /** Variable to define whether we need to update solver constants on the next iteration */ bool m_updateSolverConstants; @@ -281,7 +283,7 @@ protected: /** * Collision shapes being passed across to the cloths in this solver. */ - btAlignedObjectArray< CollisionShapeDescription > m_collisionObjectDetails; + btAlignedObjectArray< btCPUCollisionShapeDescription > m_collisionObjectDetails; void prepareCollisionConstraints(); @@ -298,7 +300,10 @@ protected: void applyForces( float solverdt ); void integrate( float solverdt ); void updateConstants( float timeStep ); - btAcceleratedSoftBodyInterface *findSoftBodyInterface( const btSoftBody* const softBody ); + int findSoftBodyIndex( const btSoftBody* const softBody ); + + /** Update the bounds of the soft body objects in the solver */ + void updateBounds(); public: @@ -306,6 +311,12 @@ public: virtual ~btCPUSoftBodySolver(); + + virtual SolverTypes getSolverType() const + { + return CPU_SOLVER; + } + virtual btSoftBodyLinkData &getLinkData(); @@ -315,16 +326,8 @@ public: - - - /** - * Add a collision object to be used by the indicated softbody. - */ - virtual void addCollisionObjectForSoftBody( int clothIdentifier, btCollisionObject *collisionObject ); - - - - + btAcceleratedSoftBodyInterface *findSoftBodyInterface( const btSoftBody* const softBody ); + const btAcceleratedSoftBodyInterface * const findSoftBodyInterface( const btSoftBody* const softBody ) const; @@ -332,15 +335,36 @@ public: virtual void updateSoftBodies( ); - virtual void optimize( btAlignedObjectArray< btSoftBody * > &softBodies ); + virtual void optimize( btAlignedObjectArray< btSoftBody * > &softBodies , bool forceUpdate=false); + + virtual void copyBackToSoftBodies(); virtual void solveConstraints( float solverdt ); virtual void predictMotion( float solverdt ); - virtual void copySoftBodyToVertexBuffer( const btSoftBody *const softBody, btVertexBufferDescriptor *vertexBuffer ); + virtual void processCollision( btSoftBody *, btCollisionObject* ); + + virtual void processCollision( btSoftBody*, btSoftBody *); + }; -#endif // #ifndef BT_ACCELERATED_SOFT_BODY_CPU_SOLVER_H +/** + * Class to manage movement of data from a solver to a given target. + * This version is the CPU to CPU generic version. + */ +class btSoftBodySolverOutputCPUtoCPU : public btSoftBodySolverOutput +{ +protected: +public: + btSoftBodySolverOutputCPUtoCPU() + { + } + + /** Output current computed vertex data to the vertex buffers for all cloths in the solver. */ + virtual void copySoftBodyToVertexBuffer( const btSoftBody * const softBody, btVertexBufferDescriptor *vertexBuffer ); +}; + +#endif // #ifndef BT_ACCELERATED_SOFT_BODY_CPU_SOLVER_H \ No newline at end of file diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/CMakeLists.txt b/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/CMakeLists.txt index 3954bc298..393d37d6a 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/CMakeLists.txt +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/CMakeLists.txt @@ -39,8 +39,11 @@ SET(BulletSoftBodyDX11Solvers_Shaders Integrate UpdatePositions UpdateNodes + ComputeBounds SolvePositions SolvePositionsSIMDBatched + SolveCollisionsAndUpdateVelocities + SolveCollisionsAndUpdateVelocitiesSIMDBatched UpdatePositionsFromVelocities ApplyForces PrepareLinks diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolverBuffer_DX11.h b/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolverBuffer_DX11.h index ff7f12ece..daf1ef899 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolverBuffer_DX11.h +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolverBuffer_DX11.h @@ -68,6 +68,9 @@ protected: buffer_desc.MiscFlags = D3D11_RESOURCE_MISC_BUFFER_STRUCTURED; buffer_desc.ByteWidth = m_CPUBuffer->size() * sizeof(ElementType); + // At a minimum the buffer must exist + if( buffer_desc.ByteWidth == 0 ) + buffer_desc.ByteWidth = sizeof(ElementType); buffer_desc.StructureByteStride = sizeof(ElementType); hr = m_d3dDevice->CreateBuffer(&buffer_desc, NULL, &m_Buffer); if( FAILED( hr ) ) @@ -82,6 +85,8 @@ protected: srvbuffer_desc.ViewDimension = D3D11_SRV_DIMENSION_BUFFER; srvbuffer_desc.Buffer.ElementWidth = m_CPUBuffer->size(); + if( srvbuffer_desc.Buffer.ElementWidth == 0 ) + srvbuffer_desc.Buffer.ElementWidth = 1; hr = m_d3dDevice->CreateShaderResourceView(m_Buffer, &srvbuffer_desc, &m_SRV); if( FAILED( hr ) ) return (hr==S_OK); @@ -93,6 +98,8 @@ protected: srvbuffer_desc.ViewDimension = D3D11_SRV_DIMENSION_BUFFER; srvbuffer_desc.Buffer.ElementWidth = m_CPUBuffer->size(); + if( srvbuffer_desc.Buffer.ElementWidth == 0 ) + srvbuffer_desc.Buffer.ElementWidth = 1; hr = m_d3dDevice->CreateShaderResourceView(m_Buffer, &srvbuffer_desc, &m_SRV); if( FAILED( hr ) ) return (hr==S_OK); @@ -104,6 +111,8 @@ protected: uavbuffer_desc.ViewDimension = D3D11_UAV_DIMENSION_BUFFER; uavbuffer_desc.Buffer.NumElements = m_CPUBuffer->size(); + if( uavbuffer_desc.Buffer.NumElements == 0 ) + uavbuffer_desc.Buffer.NumElements = 1; hr = m_d3dDevice->CreateUnorderedAccessView(m_Buffer, &uavbuffer_desc, &m_UAV); if( FAILED( hr ) ) return (hr==S_OK); @@ -173,7 +182,8 @@ public: */ bool moveToGPU() { - if( (m_CPUBuffer->size() != m_gpuSize) ) + // Reallocate if GPU size is too small + if( (m_CPUBuffer->size() > m_gpuSize ) ) m_onGPU = false; if( !m_onGPU && m_CPUBuffer->size() > 0 ) { @@ -192,16 +202,20 @@ public: } } - D3D11_BOX destRegion; - destRegion.left = 0; - destRegion.front = 0; - destRegion.top = 0; - destRegion.bottom = 1; - destRegion.back = 1; - destRegion.right = (m_CPUBuffer->size())*sizeof(ElementType); - m_d3dDeviceContext->UpdateSubresource(m_Buffer, 0, &destRegion, &((*m_CPUBuffer)[0]), 0, 0); + if( m_gpuSize > 0 ) + { + D3D11_BOX destRegion; + destRegion.left = 0; + destRegion.front = 0; + destRegion.top = 0; + destRegion.bottom = 1; + destRegion.back = 1; + destRegion.right = (m_CPUBuffer->size())*sizeof(ElementType); + m_d3dDeviceContext->UpdateSubresource(m_Buffer, 0, &destRegion, &((*m_CPUBuffer)[0]), 0, 0); + + m_onGPU = true; + } - m_onGPU = true; } return true; diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11.cpp b/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11.cpp index 7877aa6a0..b0227e468 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11.cpp +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11.cpp @@ -19,6 +19,7 @@ subject to the following restrictions: #include "btSoftBodySolver_DX11.h" #include "btSoftBodySolverVertexBuffer_DX11.h" #include "BulletSoftBody/btSoftBody.h" +#include "BulletCollision/CollisionShapes/btCapsuleShape.h" #define MSTRINGIFY(A) #A static char* PrepareLinksHLSLString = @@ -43,6 +44,10 @@ static char* OutputToVertexArrayHLSLString = #include "HLSL/OutputToVertexArray.hlsl" static char* VSolveLinksHLSLString = #include "HLSL/VSolveLinks.hlsl" +static char* ComputeBoundsHLSLString = +#include "HLSL/ComputeBounds.hlsl" +static char* SolveCollisionsAndUpdateVelocitiesHLSLString = +#include "HLSL/SolveCollisionsAndUpdateVelocities.hlsl" btSoftBodyLinkDataDX11::btSoftBodyLinkDataDX11( ID3D11Device *d3dDevice, ID3D11DeviceContext *d3dDeviceContext ) : @@ -545,6 +550,7 @@ void btSoftBodyTriangleDataDX11::generateBatches() btDX11SoftBodySolver::btDX11SoftBodySolver(ID3D11Device * dx11Device, ID3D11DeviceContext* dx11Context) : m_dx11Device( dx11Device ), m_dx11Context( dx11Context ), + dxFunctions( m_dx11Device, m_dx11Context ), m_linkData(m_dx11Device, m_dx11Context), m_vertexData(m_dx11Device, m_dx11Context), m_triangleData(m_dx11Device, m_dx11Context), @@ -554,7 +560,12 @@ btDX11SoftBodySolver::btDX11SoftBodySolver(ID3D11Device * dx11Device, ID3D11Devi m_dx11PerClothVelocityCorrectionCoefficient( m_dx11Device, m_dx11Context, &m_perClothVelocityCorrectionCoefficient, true ), m_dx11PerClothLiftFactor( m_dx11Device, m_dx11Context, &m_perClothLiftFactor, true ), m_dx11PerClothDragFactor( m_dx11Device, m_dx11Context, &m_perClothDragFactor, true ), - m_dx11PerClothMediumDensity( m_dx11Device, m_dx11Context, &m_perClothMediumDensity, true ) + m_dx11PerClothMediumDensity( m_dx11Device, m_dx11Context, &m_perClothMediumDensity, true ), + m_dx11PerClothCollisionObjects( m_dx11Device, m_dx11Context, &m_perClothCollisionObjects, true ), + m_dx11CollisionObjectDetails( m_dx11Device, m_dx11Context, &m_collisionObjectDetails, true ), + m_dx11PerClothMinBounds( m_dx11Device, m_dx11Context, &m_perClothMinBounds, false ), + m_dx11PerClothMaxBounds( m_dx11Device, m_dx11Context, &m_perClothMaxBounds, false ), + m_dx11PerClothFriction( m_dx11Device, m_dx11Context, &m_perClothFriction, false ) { // Initial we will clearly need to update solver constants // For now this is global for the cloths linked with this solver - we should probably make this body specific @@ -565,15 +576,20 @@ btDX11SoftBodySolver::btDX11SoftBodySolver(ID3D11Device * dx11Device, ID3D11Devi } btDX11SoftBodySolver::~btDX11SoftBodySolver() +{ + releaseKernels(); +} + +void btDX11SoftBodySolver::releaseKernels() { + + SAFE_RELEASE( prepareLinksKernel.kernel ); + SAFE_RELEASE( prepareLinksKernel.constBuffer ); + SAFE_RELEASE( integrateKernel.kernel ); SAFE_RELEASE( integrateKernel.constBuffer ); SAFE_RELEASE( integrateKernel.kernel ); - SAFE_RELEASE( prepareLinksKernel.constBuffer ); - SAFE_RELEASE( prepareLinksKernel.kernel ); SAFE_RELEASE( solvePositionsFromLinksKernel.constBuffer ); SAFE_RELEASE( solvePositionsFromLinksKernel.kernel ); - SAFE_RELEASE( vSolveLinksKernel.constBuffer ); - SAFE_RELEASE( vSolveLinksKernel.kernel ); SAFE_RELEASE( updatePositionsFromVelocitiesKernel.constBuffer ); SAFE_RELEASE( updatePositionsFromVelocitiesKernel.kernel ); SAFE_RELEASE( updateVelocitiesFromPositionsWithoutVelocitiesKernel.constBuffer ); @@ -586,27 +602,57 @@ btDX11SoftBodySolver::~btDX11SoftBodySolver() SAFE_RELEASE( normalizeNormalsAndAreasKernel.kernel ); SAFE_RELEASE( updateSoftBodiesKernel.constBuffer ); SAFE_RELEASE( updateSoftBodiesKernel.kernel ); - SAFE_RELEASE( outputToVertexArrayWithNormalsKernel.constBuffer ); - SAFE_RELEASE( outputToVertexArrayWithNormalsKernel.kernel ); - SAFE_RELEASE( outputToVertexArrayWithoutNormalsKernel.constBuffer ); - SAFE_RELEASE( outputToVertexArrayWithoutNormalsKernel.kernel ); - + SAFE_RELEASE( solveCollisionsAndUpdateVelocitiesKernel.kernel ); + SAFE_RELEASE( solveCollisionsAndUpdateVelocitiesKernel.constBuffer ); + SAFE_RELEASE( computeBoundsKernel.kernel ); + SAFE_RELEASE( computeBoundsKernel.constBuffer ); + SAFE_RELEASE( vSolveLinksKernel.kernel ); + SAFE_RELEASE( vSolveLinksKernel.constBuffer ); SAFE_RELEASE( addVelocityKernel.constBuffer ); SAFE_RELEASE( addVelocityKernel.kernel ); SAFE_RELEASE( applyForcesKernel.constBuffer ); SAFE_RELEASE( applyForcesKernel.kernel ); - SAFE_RELEASE( outputToVertexArrayKernel.constBuffer ); - SAFE_RELEASE( outputToVertexArrayKernel.kernel ); - SAFE_RELEASE( collideCylinderKernel.constBuffer ); - SAFE_RELEASE( collideCylinderKernel.kernel ); + + m_shadersInitialized = false; } - -void btDX11SoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * > &softBodies ) +void btDX11SoftBodySolver::copyBackToSoftBodies() { - if( m_softBodySet.size() != softBodies.size() ) + // Move the vertex data back to the host first + m_vertexData.moveFromAccelerator(); + + // Loop over soft bodies, copying all the vertex positions back for each body in turn + for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) + { + btAcceleratedSoftBodyInterface *softBodyInterface = m_softBodySet[ softBodyIndex ]; + btSoftBody *softBody = softBodyInterface->getSoftBody(); + + int firstVertex = softBodyInterface->getFirstVertex(); + int numVertices = softBodyInterface->getNumVertices(); + + // Copy vertices from solver back into the softbody + for( int vertex = 0; vertex < numVertices; ++vertex ) + { + using Vectormath::Aos::Point3; + Point3 vertexPosition( getVertexData().getVertexPositions()[firstVertex + vertex] ); + + softBody->m_nodes[vertex].m_x.setX( vertexPosition.getX() ); + softBody->m_nodes[vertex].m_x.setY( vertexPosition.getY() ); + softBody->m_nodes[vertex].m_x.setZ( vertexPosition.getZ() ); + + softBody->m_nodes[vertex].m_n.setX( vertexPosition.getX() ); + softBody->m_nodes[vertex].m_n.setY( vertexPosition.getY() ); + softBody->m_nodes[vertex].m_n.setZ( vertexPosition.getZ() ); + } + } +} // btDX11SoftBodySolver::copyBackToSoftBodies + + +void btDX11SoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * > &softBodies, bool forceUpdate ) +{ + if( forceUpdate || m_softBodySet.size() != softBodies.size() ) { // Have a change in the soft body set so update, reloading all the data getVertexData().clear(); @@ -622,7 +668,7 @@ void btDX11SoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * > &softB using Vectormath::Aos::Point3; // Create SoftBody that will store the information within the solver - btDX11AcceleratedSoftBodyInterface *newSoftBody = new btDX11AcceleratedSoftBodyInterface( softBody ); + btAcceleratedSoftBodyInterface *newSoftBody = new btAcceleratedSoftBodyInterface( softBody ); m_softBodySet.push_back( newSoftBody ); m_perClothAcceleration.push_back( toVector3(softBody->getWorldInfo()->m_gravity) ); @@ -631,6 +677,11 @@ void btDX11SoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * > &softB m_perClothLiftFactor.push_back( softBody->m_cfg.kLF ); m_perClothDragFactor.push_back( softBody->m_cfg.kDG ); m_perClothMediumDensity.push_back(softBody->getWorldInfo()->air_density); + // Simple init values. Actually we'll put 0 and -1 into them at the appropriate time + m_perClothMinBounds.push_back( UIntVector3( 0, 0, 0 ) ); + m_perClothMaxBounds.push_back( UIntVector3( UINT_MAX, UINT_MAX, UINT_MAX ) ); + m_perClothFriction.push_back( softBody->getFriction() ); + m_perClothCollisionObjects.push_back( CollisionObjectIndices(-1, -1) ); // Add space for new vertices and triangles in the default solver for now // TODO: Include space here for tearing too later @@ -738,7 +789,11 @@ btSoftBodyTriangleData &btDX11SoftBodySolver::getTriangleData() bool btDX11SoftBodySolver::checkInitialized() { - return buildShaders(); + if( !m_shadersInitialized ) + if( buildShaders() ) + m_shadersInitialized = true; + + return m_shadersInitialized; } void btDX11SoftBodySolver::resetNormalsAndAreas( int numVertices ) @@ -892,6 +947,7 @@ void btDX11SoftBodySolver::updateSoftBodies() normalizeNormalsAndAreas( numVertices ); + } // btDX11SoftBodySolver::updateSoftBodies @@ -1046,6 +1102,80 @@ float btDX11SoftBodySolver::computeTriangleArea( return area; } // btDX11SoftBodySolver::computeTriangleArea + +void btDX11SoftBodySolver::updateBounds() +{ + using Vectormath::Aos::Point3; + // Interpretation structure for float and int + + struct FPRep { + unsigned int mantissa : 23; + unsigned int exponent : 8; + unsigned int sign : 1; + }; + union FloatAsInt + { + float floatValue; + int intValue; + unsigned int uintValue; + FPRep fpRep; + }; + + + // Update bounds array to min and max int values to allow easy atomics + for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) + { + m_perClothMinBounds[softBodyIndex] = UIntVector3( UINT_MAX, UINT_MAX, UINT_MAX ); + m_perClothMaxBounds[softBodyIndex] = UIntVector3( 0, 0, 0 ); + } + + m_dx11PerClothMinBounds.moveToGPU(); + m_dx11PerClothMaxBounds.moveToGPU(); + + + computeBounds( ); + + + m_dx11PerClothMinBounds.moveFromGPU(); + m_dx11PerClothMaxBounds.moveFromGPU(); + + + + for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) + { + UIntVector3 minBoundUInt = m_perClothMinBounds[softBodyIndex]; + UIntVector3 maxBoundUInt = m_perClothMaxBounds[softBodyIndex]; + + // Convert back to float + FloatAsInt fai; + + btVector3 minBound; + fai.uintValue = minBoundUInt.x; + fai.uintValue ^= (((fai.uintValue >> 31) - 1) | 0x80000000); + minBound.setX( fai.floatValue ); + fai.uintValue = minBoundUInt.y; + fai.uintValue ^= (((fai.uintValue >> 31) - 1) | 0x80000000); + minBound.setY( fai.floatValue ); + fai.uintValue = minBoundUInt.z; + fai.uintValue ^= (((fai.uintValue >> 31) - 1) | 0x80000000); + minBound.setZ( fai.floatValue ); + + btVector3 maxBound; + fai.uintValue = maxBoundUInt.x; + fai.uintValue ^= (((fai.uintValue >> 31) - 1) | 0x80000000); + maxBound.setX( fai.floatValue ); + fai.uintValue = maxBoundUInt.y; + fai.uintValue ^= (((fai.uintValue >> 31) - 1) | 0x80000000); + maxBound.setY( fai.floatValue ); + fai.uintValue = maxBoundUInt.z; + fai.uintValue ^= (((fai.uintValue >> 31) - 1) | 0x80000000); + maxBound.setZ( fai.floatValue ); + + // And finally assign to the soft body + m_softBodySet[softBodyIndex]->updateBounds( minBound, maxBound ); + } +} + void btDX11SoftBodySolver::updateConstants( float timeStep ) { using namespace Vectormath::Aos; @@ -1074,6 +1204,59 @@ void btDX11SoftBodySolver::updateConstants( float timeStep ) } } // btDX11SoftBodySolver::updateConstants +/** + * Sort the collision object details array and generate indexing into it for the per-cloth collision object array. + */ +void btDX11SoftBodySolver::prepareCollisionConstraints() +{ + // First do a simple sort on the collision objects + btAlignedObjectArray numObjectsPerClothPrefixSum; + btAlignedObjectArray numObjectsPerCloth; + numObjectsPerCloth.resize( m_softBodySet.size(), 0 ); + numObjectsPerClothPrefixSum.resize( m_softBodySet.size(), 0 ); + + + class QuickSortCompare + { + public: + + bool operator() ( const CollisionShapeDescription& a, const CollisionShapeDescription& b ) + { + return ( a.softBodyIdentifier < b.softBodyIdentifier ); + } + }; + + QuickSortCompare comparator; + m_collisionObjectDetails.quickSort( comparator ); + + // Generating indexing for perClothCollisionObjects + // First clear the previous values with the "no collision object for cloth" constant + for( int clothIndex = 0; clothIndex < m_perClothCollisionObjects.size(); ++clothIndex ) + { + m_perClothCollisionObjects[clothIndex].firstObject = -1; + m_perClothCollisionObjects[clothIndex].endObject = -1; + } + int currentCloth = 0; + int startIndex = 0; + for( int collisionObject = 0; collisionObject < m_collisionObjectDetails.size(); ++collisionObject ) + { + int nextCloth = m_collisionObjectDetails[collisionObject].softBodyIdentifier; + if( nextCloth != currentCloth ) + { + // Changed cloth in the array + // Set the end index and the range is what we need for currentCloth + m_perClothCollisionObjects[currentCloth].firstObject = startIndex; + m_perClothCollisionObjects[currentCloth].endObject = collisionObject; + currentCloth = nextCloth; + startIndex = collisionObject; + } + } + + // And update last cloth + m_perClothCollisionObjects[currentCloth].firstObject = startIndex; + m_perClothCollisionObjects[currentCloth].endObject = m_collisionObjectDetails.size(); + +} // btDX11SoftBodySolver::prepareCollisionConstraints void btDX11SoftBodySolver::solveConstraints( float solverdt ) @@ -1115,6 +1298,9 @@ void btDX11SoftBodySolver::solveConstraints( float solverdt ) } } + + prepareCollisionConstraints(); + // Compute new positions from velocity // Also update the previous position so that our position computation is now based on the new position from the velocity solution // rather than based directly on the original positions @@ -1139,7 +1325,8 @@ void btDX11SoftBodySolver::solveConstraints( float solverdt ) } // for( int iteration = 0; iteration < m_numberOfPositionIterations ; ++iteration ) - updateVelocitiesFromPositionsWithoutVelocities( 1.f/solverdt ); + // At this point assume that the force array is blank - we will overwrite it + solveCollisionsAndUpdateVelocities( 1.f/solverdt ); } // btDX11SoftBodySolver::solveConstraints @@ -1435,6 +1622,114 @@ void btDX11SoftBodySolver::updateVelocitiesFromPositionsWithoutVelocities( float } // btDX11SoftBodySolver::updateVelocitiesFromPositionsWithoutVelocities + +void btDX11SoftBodySolver::computeBounds( ) +{ + ComputeBoundsCB constBuffer; + m_vertexData.moveToAccelerator(); + + // Set the first link of the batch + // and the batch size + constBuffer.numNodes = m_vertexData.getNumVertices(); + constBuffer.numSoftBodies = m_softBodySet.size(); + + D3D11_MAPPED_SUBRESOURCE MappedResource = {0}; + m_dx11Context->Map( computeBoundsKernel.constBuffer, 0, D3D11_MAP_WRITE_DISCARD, 0, &MappedResource ); + memcpy( MappedResource.pData, &constBuffer, sizeof(ComputeBoundsCB) ); + m_dx11Context->Unmap( computeBoundsKernel.constBuffer, 0 ); + m_dx11Context->CSSetConstantBuffers( 0, 1, &computeBoundsKernel.constBuffer ); + + // Set resources and dispatch + m_dx11Context->CSSetShaderResources( 0, 1, &(m_vertexData.m_dx11ClothIdentifier.getSRV()) ); + m_dx11Context->CSSetShaderResources( 1, 1, &(m_vertexData.m_dx11VertexPosition.getSRV()) ); + + m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &(m_dx11PerClothMinBounds.getUAV()), NULL ); + m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &(m_dx11PerClothMaxBounds.getUAV()), NULL ); + + // Execute the kernel + m_dx11Context->CSSetShader( computeBoundsKernel.kernel, NULL, 0 ); + + int numBlocks = (constBuffer.numNodes + (128-1)) / 128; + m_dx11Context->Dispatch(numBlocks , 1, 1 ); + + { + // Tidy up + ID3D11ShaderResourceView* pViewNULL = NULL; + m_dx11Context->CSSetShaderResources( 0, 1, &pViewNULL ); + m_dx11Context->CSSetShaderResources( 1, 1, &pViewNULL ); + + ID3D11UnorderedAccessView* pUAViewNULL = NULL; + m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &pUAViewNULL, NULL ); + m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &pUAViewNULL, NULL ); + + ID3D11Buffer *pBufferNull = NULL; + m_dx11Context->CSSetConstantBuffers( 0, 1, &pBufferNull ); + } +} + +void btDX11SoftBodySolver::solveCollisionsAndUpdateVelocities( float isolverdt ) +{ + + // Copy kernel parameters to GPU + m_vertexData.moveToAccelerator(); + m_dx11PerClothFriction.moveToGPU(); + m_dx11PerClothDampingFactor.moveToGPU(); + m_dx11PerClothCollisionObjects.moveToGPU(); + m_dx11CollisionObjectDetails.moveToGPU(); + + SolveCollisionsAndUpdateVelocitiesCB constBuffer; + + // Set the first link of the batch + // and the batch size + constBuffer.numNodes = m_vertexData.getNumVertices(); + constBuffer.isolverdt = isolverdt; + + + D3D11_MAPPED_SUBRESOURCE MappedResource = {0}; + m_dx11Context->Map( solveCollisionsAndUpdateVelocitiesKernel.constBuffer, 0, D3D11_MAP_WRITE_DISCARD, 0, &MappedResource ); + memcpy( MappedResource.pData, &constBuffer, sizeof(SolveCollisionsAndUpdateVelocitiesCB) ); + m_dx11Context->Unmap( solveCollisionsAndUpdateVelocitiesKernel.constBuffer, 0 ); + m_dx11Context->CSSetConstantBuffers( 0, 1, &solveCollisionsAndUpdateVelocitiesKernel.constBuffer ); + + // Set resources and dispatch + m_dx11Context->CSSetShaderResources( 0, 1, &(m_vertexData.m_dx11ClothIdentifier.getSRV()) ); + m_dx11Context->CSSetShaderResources( 1, 1, &(m_vertexData.m_dx11VertexPreviousPosition.getSRV()) ); + m_dx11Context->CSSetShaderResources( 2, 1, &(m_dx11PerClothFriction.getSRV()) ); + m_dx11Context->CSSetShaderResources( 3, 1, &(m_dx11PerClothDampingFactor.getSRV()) ); + m_dx11Context->CSSetShaderResources( 4, 1, &(m_dx11PerClothCollisionObjects.getSRV()) ); + m_dx11Context->CSSetShaderResources( 5, 1, &(m_dx11CollisionObjectDetails.getSRV()) ); + + m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &(m_vertexData.m_dx11VertexForceAccumulator.getUAV()), NULL ); + m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &(m_vertexData.m_dx11VertexVelocity.getUAV()), NULL ); + m_dx11Context->CSSetUnorderedAccessViews( 2, 1, &(m_vertexData.m_dx11VertexPosition.getUAV()), NULL ); + + // Execute the kernel + m_dx11Context->CSSetShader( solveCollisionsAndUpdateVelocitiesKernel.kernel, NULL, 0 ); + + int numBlocks = (constBuffer.numNodes + (128-1)) / 128; + m_dx11Context->Dispatch(numBlocks , 1, 1 ); + + { + // Tidy up + ID3D11ShaderResourceView* pViewNULL = NULL; + m_dx11Context->CSSetShaderResources( 0, 1, &pViewNULL ); + m_dx11Context->CSSetShaderResources( 1, 1, &pViewNULL ); + m_dx11Context->CSSetShaderResources( 2, 1, &pViewNULL ); + m_dx11Context->CSSetShaderResources( 3, 1, &pViewNULL ); + m_dx11Context->CSSetShaderResources( 4, 1, &pViewNULL ); + m_dx11Context->CSSetShaderResources( 5, 1, &pViewNULL ); + + ID3D11UnorderedAccessView* pUAViewNULL = NULL; + m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &pUAViewNULL, NULL ); + m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &pUAViewNULL, NULL ); + m_dx11Context->CSSetUnorderedAccessViews( 2, 1, &pUAViewNULL, NULL ); + + ID3D11Buffer *pBufferNull = NULL; + m_dx11Context->CSSetConstantBuffers( 0, 1, &pBufferNull ); + } + +} // btDX11SoftBodySolver::solveCollisionsAndUpdateVelocities + // End kernel dispatches ///////////////////////////////////// @@ -1451,23 +1746,51 @@ void btDX11SoftBodySolver::updateVelocitiesFromPositionsWithoutVelocities( float -btDX11AcceleratedSoftBodyInterface *btDX11SoftBodySolver::findSoftBodyInterface( const btSoftBody* const softBody ) +btDX11SoftBodySolver::btAcceleratedSoftBodyInterface *btDX11SoftBodySolver::findSoftBodyInterface( const btSoftBody* const softBody ) { for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) { - btDX11AcceleratedSoftBodyInterface *softBodyInterface = m_softBodySet[softBodyIndex]; + btAcceleratedSoftBodyInterface *softBodyInterface = m_softBodySet[softBodyIndex]; if( softBodyInterface->getSoftBody() == softBody ) return softBodyInterface; } return 0; } -void btDX11SoftBodySolver::copySoftBodyToVertexBuffer( const btSoftBody * const softBody, btVertexBufferDescriptor *vertexBuffer ) +const btDX11SoftBodySolver::btAcceleratedSoftBodyInterface * const btDX11SoftBodySolver::findSoftBodyInterface( const btSoftBody* const softBody ) const { - checkInitialized(); - - btDX11AcceleratedSoftBodyInterface *currentCloth = findSoftBodyInterface( softBody ); + for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) + { + btAcceleratedSoftBodyInterface *softBodyInterface = m_softBodySet[softBodyIndex]; + if( softBodyInterface->getSoftBody() == softBody ) + return softBodyInterface; + } + return 0; +} +int btDX11SoftBodySolver::findSoftBodyIndex( const btSoftBody* const softBody ) +{ + for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) + { + btAcceleratedSoftBodyInterface *softBodyInterface = m_softBodySet[softBodyIndex]; + if( softBodyInterface->getSoftBody() == softBody ) + return softBodyIndex; + } + return 1; +} + + +void btSoftBodySolverOutputDXtoCPU::copySoftBodyToVertexBuffer( const btSoftBody * const softBody, btVertexBufferDescriptor *vertexBuffer ) +{ + + + btSoftBodySolver *solver = softBody->getSoftBodySolver(); + btAssert( solver->getSolverType() == btSoftBodySolver::DX_SOLVER || solver->getSolverType() == btSoftBodySolver::DX_SIMD_SOLVER ); + btDX11SoftBodySolver *dxSolver = static_cast< btDX11SoftBodySolver * >( solver ); + + btDX11SoftBodySolver::btAcceleratedSoftBodyInterface * currentCloth = dxSolver->findSoftBodyInterface( softBody ); + btSoftBodyVertexDataDX11 &vertexData( dxSolver->m_vertexData ); + const int firstVertex = currentCloth->getFirstVertex(); const int lastVertex = firstVertex + currentCloth->getNumVertices(); @@ -1475,8 +1798,8 @@ void btDX11SoftBodySolver::copySoftBodyToVertexBuffer( const btSoftBody * const if( vertexBuffer->getBufferType() == btVertexBufferDescriptor::CPU_BUFFER ) { // If we're doing a CPU-buffer copy must copy the data back to the host first - m_vertexData.m_dx11VertexPosition.copyFromGPU(); - m_vertexData.m_dx11VertexNormal.copyFromGPU(); + vertexData.m_dx11VertexPosition.copyFromGPU(); + vertexData.m_dx11VertexNormal.copyFromGPU(); const int firstVertex = currentCloth->getFirstVertex(); const int lastVertex = firstVertex + currentCloth->getNumVertices(); @@ -1491,7 +1814,7 @@ void btDX11SoftBodySolver::copySoftBodyToVertexBuffer( const btSoftBody * const for( int vertexIndex = firstVertex; vertexIndex < lastVertex; ++vertexIndex ) { - Vectormath::Aos::Point3 position = m_vertexData.getPosition(vertexIndex); + Vectormath::Aos::Point3 position = vertexData.getPosition(vertexIndex); *(vertexPointer + 0) = position.getX(); *(vertexPointer + 1) = position.getY(); *(vertexPointer + 2) = position.getZ(); @@ -1506,13 +1829,82 @@ void btDX11SoftBodySolver::copySoftBodyToVertexBuffer( const btSoftBody * const for( int vertexIndex = firstVertex; vertexIndex < lastVertex; ++vertexIndex ) { - Vectormath::Aos::Vector3 normal = m_vertexData.getNormal(vertexIndex); + Vectormath::Aos::Vector3 normal = vertexData.getNormal(vertexIndex); *(normalPointer + 0) = normal.getX(); *(normalPointer + 1) = normal.getY(); *(normalPointer + 2) = normal.getZ(); normalPointer += normalStride; } } + } +} // btDX11SoftBodySolver::outputToVertexBuffers + + + +bool btSoftBodySolverOutputDXtoDX::checkInitialized() +{ + if( !m_shadersInitialized ) + if( buildShaders() ) + m_shadersInitialized = true; + + return m_shadersInitialized; +} + +void btSoftBodySolverOutputDXtoDX::releaseKernels() +{ + SAFE_RELEASE( outputToVertexArrayWithNormalsKernel.constBuffer ); + SAFE_RELEASE( outputToVertexArrayWithNormalsKernel.kernel ); + SAFE_RELEASE( outputToVertexArrayWithoutNormalsKernel.constBuffer ); + SAFE_RELEASE( outputToVertexArrayWithoutNormalsKernel.kernel ); + + m_shadersInitialized = false; +} + + +bool btSoftBodySolverOutputDXtoDX::buildShaders() +{ + // Ensure current kernels are released first + releaseKernels(); + + bool returnVal = true; + + if( m_shadersInitialized ) + return true; + + + outputToVertexArrayWithNormalsKernel = dxFunctions.compileComputeShaderFromString( OutputToVertexArrayHLSLString, "OutputToVertexArrayWithNormalsKernel", sizeof(OutputToVertexArrayCB) ); + if( !outputToVertexArrayWithNormalsKernel.constBuffer) + returnVal = false; + outputToVertexArrayWithoutNormalsKernel = dxFunctions.compileComputeShaderFromString( OutputToVertexArrayHLSLString, "OutputToVertexArrayWithoutNormalsKernel", sizeof(OutputToVertexArrayCB) ); + if( !outputToVertexArrayWithoutNormalsKernel.constBuffer ) + returnVal = false; + + + if( returnVal ) + m_shadersInitialized = true; + + return returnVal; +} + + +void btSoftBodySolverOutputDXtoDX::copySoftBodyToVertexBuffer( const btSoftBody * const softBody, btVertexBufferDescriptor *vertexBuffer ) +{ + + + btSoftBodySolver *solver = softBody->getSoftBodySolver(); + btAssert( solver->getSolverType() == btSoftBodySolver::DX_SOLVER || solver->getSolverType() == btSoftBodySolver::DX_SIMD_SOLVER ); + btDX11SoftBodySolver *dxSolver = static_cast< btDX11SoftBodySolver * >( solver ); + checkInitialized(); + btDX11SoftBodySolver::btAcceleratedSoftBodyInterface * currentCloth = dxSolver->findSoftBodyInterface( softBody ); + btSoftBodyVertexDataDX11 &vertexData( dxSolver->m_vertexData ); + + + const int firstVertex = currentCloth->getFirstVertex(); + const int lastVertex = firstVertex + currentCloth->getNumVertices(); + + if( vertexBuffer->getBufferType() == btVertexBufferDescriptor::CPU_BUFFER ) + { + btSoftBodySolverOutputDXtoDX::copySoftBodyToVertexBuffer( softBody, vertexBuffer ); } else if( vertexBuffer->getBufferType() == btVertexBufferDescriptor::DX11_BUFFER ) { // Do a DX11 copy shader DX to DX copy @@ -1539,86 +1931,43 @@ void btDX11SoftBodySolver::copySoftBodyToVertexBuffer( const btSoftBody * const // TODO: factor this out. Number of nodes is static and sdt might be, too, we can update this just once on setup D3D11_MAPPED_SUBRESOURCE MappedResource = {0}; - m_dx11Context->Map( outputToVertexArrayConstBuffer, 0, D3D11_MAP_WRITE_DISCARD, 0, &MappedResource ); + dxFunctions.m_dx11Context->Map( outputToVertexArrayConstBuffer, 0, D3D11_MAP_WRITE_DISCARD, 0, &MappedResource ); memcpy( MappedResource.pData, &constBuffer, sizeof(OutputToVertexArrayCB) ); - m_dx11Context->Unmap( outputToVertexArrayConstBuffer, 0 ); - m_dx11Context->CSSetConstantBuffers( 0, 1, &outputToVertexArrayConstBuffer ); + dxFunctions.m_dx11Context->Unmap( outputToVertexArrayConstBuffer, 0 ); + dxFunctions.m_dx11Context->CSSetConstantBuffers( 0, 1, &outputToVertexArrayConstBuffer ); // Set resources and dispatch - m_dx11Context->CSSetShaderResources( 0, 1, &(m_vertexData.m_dx11VertexPosition.getSRV()) ); - m_dx11Context->CSSetShaderResources( 1, 1, &(m_vertexData.m_dx11VertexNormal.getSRV()) ); + dxFunctions.m_dx11Context->CSSetShaderResources( 0, 1, &(vertexData.m_dx11VertexPosition.getSRV()) ); + dxFunctions.m_dx11Context->CSSetShaderResources( 1, 1, &(vertexData.m_dx11VertexNormal.getSRV()) ); ID3D11UnorderedAccessView* dx11UAV = dx11VertexBuffer->getDX11UAV(); - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &(dx11UAV), NULL ); + dxFunctions.m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &(dx11UAV), NULL ); // Execute the kernel - m_dx11Context->CSSetShader( outputToVertexArrayShader, NULL, 0 ); + dxFunctions.m_dx11Context->CSSetShader( outputToVertexArrayShader, NULL, 0 ); int numBlocks = (constBuffer.numNodes + (128-1)) / 128; - m_dx11Context->Dispatch(numBlocks, 1, 1 ); + dxFunctions.m_dx11Context->Dispatch(numBlocks, 1, 1 ); { // Tidy up ID3D11ShaderResourceView* pViewNULL = NULL; - m_dx11Context->CSSetShaderResources( 0, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 1, 1, &pViewNULL ); + dxFunctions.m_dx11Context->CSSetShaderResources( 0, 1, &pViewNULL ); + dxFunctions.m_dx11Context->CSSetShaderResources( 1, 1, &pViewNULL ); ID3D11UnorderedAccessView* pUAViewNULL = NULL; - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &pUAViewNULL, NULL ); + dxFunctions.m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &pUAViewNULL, NULL ); ID3D11Buffer *pBufferNull = NULL; - m_dx11Context->CSSetConstantBuffers( 0, 1, &pBufferNull ); + dxFunctions.m_dx11Context->CSSetConstantBuffers( 0, 1, &pBufferNull ); } } - - - - - - if( vertexBuffer->getBufferType() == btVertexBufferDescriptor::CPU_BUFFER ) - { - const int firstVertex = currentCloth->getFirstVertex(); - const int lastVertex = firstVertex + currentCloth->getNumVertices(); - const btCPUVertexBufferDescriptor *cpuVertexBuffer = static_cast< btCPUVertexBufferDescriptor* >(vertexBuffer); - float *basePointer = cpuVertexBuffer->getBasePointer(); - - if( vertexBuffer->hasVertexPositions() ) - { - const int vertexOffset = cpuVertexBuffer->getVertexOffset(); - const int vertexStride = cpuVertexBuffer->getVertexStride(); - float *vertexPointer = basePointer + vertexOffset; - - for( int vertexIndex = firstVertex; vertexIndex < lastVertex; ++vertexIndex ) - { - Vectormath::Aos::Point3 position = m_vertexData.getPosition(vertexIndex); - *(vertexPointer + 0) = position.getX(); - *(vertexPointer + 1) = position.getY(); - *(vertexPointer + 2) = position.getZ(); - vertexPointer += vertexStride; - } - } - if( vertexBuffer->hasNormals() ) - { - const int normalOffset = cpuVertexBuffer->getNormalOffset(); - const int normalStride = cpuVertexBuffer->getNormalStride(); - float *normalPointer = basePointer + normalOffset; - - for( int vertexIndex = firstVertex; vertexIndex < lastVertex; ++vertexIndex ) - { - Vectormath::Aos::Vector3 normal = m_vertexData.getNormal(vertexIndex); - *(normalPointer + 0) = normal.getX(); - *(normalPointer + 1) = normal.getY(); - *(normalPointer + 2) = normal.getZ(); - normalPointer += normalStride; - } - } - } } // btDX11SoftBodySolver::outputToVertexBuffers -btDX11SoftBodySolver::KernelDesc btDX11SoftBodySolver::compileComputeShaderFromString( const char* shaderString, const char* shaderName, int constBufferSize ) +DXFunctions::KernelDesc DXFunctions::compileComputeShaderFromString( const char* shaderString, const char* shaderName, int constBufferSize, D3D10_SHADER_MACRO *compileMacros ) { const char *cs5String = "cs_5_0"; @@ -1626,12 +1975,12 @@ btDX11SoftBodySolver::KernelDesc btDX11SoftBodySolver::compileComputeShaderFromS ID3DBlob* pErrorBlob = NULL; ID3DBlob* pBlob = NULL; ID3D11ComputeShader* kernelPointer = 0; - + hr = D3DX11CompileFromMemory( shaderString, strlen(shaderString), - shaderName, // file name - NULL, + shaderName, + compileMacros, NULL, shaderName, cs5String, @@ -1647,13 +1996,14 @@ btDX11SoftBodySolver::KernelDesc btDX11SoftBodySolver::compileComputeShaderFromS { if( pErrorBlob ) { btAssert( "Compilation of compute shader failed\n" ); - //OutputDebugStringA( (char*)pErrorBlob->GetBufferPointer() ); + char *debugString = (char*)pErrorBlob->GetBufferPointer(); + OutputDebugStringA( debugString ); } SAFE_RELEASE( pErrorBlob ); SAFE_RELEASE( pBlob ); - btDX11SoftBodySolver::KernelDesc descriptor; + DXFunctions::KernelDesc descriptor; descriptor.kernel = 0; descriptor.constBuffer = 0; return descriptor; @@ -1663,7 +2013,7 @@ btDX11SoftBodySolver::KernelDesc btDX11SoftBodySolver::compileComputeShaderFromS hr = m_dx11Device->CreateComputeShader( pBlob->GetBufferPointer(), pBlob->GetBufferSize(), NULL, &kernelPointer ); if( FAILED( hr ) ) { - btDX11SoftBodySolver::KernelDesc descriptor; + DXFunctions::KernelDesc descriptor; descriptor.kernel = 0; descriptor.constBuffer = 0; return descriptor; @@ -1692,60 +2042,65 @@ btDX11SoftBodySolver::KernelDesc btDX11SoftBodySolver::compileComputeShaderFromS SAFE_RELEASE( pErrorBlob ); SAFE_RELEASE( pBlob ); - btDX11SoftBodySolver::KernelDesc descriptor; + DXFunctions::KernelDesc descriptor; descriptor.kernel = kernelPointer; descriptor.constBuffer = constBuffer; return descriptor; } // compileComputeShader + bool btDX11SoftBodySolver::buildShaders() { + // Ensure current kernels are released first + releaseKernels(); + bool returnVal = true; if( m_shadersInitialized ) return true; - prepareLinksKernel = compileComputeShaderFromString( PrepareLinksHLSLString, "PrepareLinksKernel", sizeof(PrepareLinksCB) ); + prepareLinksKernel = dxFunctions.compileComputeShaderFromString( PrepareLinksHLSLString, "PrepareLinksKernel", sizeof(PrepareLinksCB) ); if( !prepareLinksKernel.constBuffer ) returnVal = false; - updatePositionsFromVelocitiesKernel = compileComputeShaderFromString( UpdatePositionsFromVelocitiesHLSLString, "UpdatePositionsFromVelocitiesKernel", sizeof(UpdatePositionsFromVelocitiesCB) ); + updatePositionsFromVelocitiesKernel = dxFunctions.compileComputeShaderFromString( UpdatePositionsFromVelocitiesHLSLString, "UpdatePositionsFromVelocitiesKernel", sizeof(UpdatePositionsFromVelocitiesCB) ); if( !updatePositionsFromVelocitiesKernel.constBuffer ) returnVal = false; - solvePositionsFromLinksKernel = compileComputeShaderFromString( SolvePositionsHLSLString, "SolvePositionsFromLinksKernel", sizeof(SolvePositionsFromLinksKernelCB) ); + solvePositionsFromLinksKernel = dxFunctions.compileComputeShaderFromString( SolvePositionsHLSLString, "SolvePositionsFromLinksKernel", sizeof(SolvePositionsFromLinksKernelCB) ); if( !updatePositionsFromVelocitiesKernel.constBuffer ) returnVal = false; - vSolveLinksKernel = compileComputeShaderFromString( VSolveLinksHLSLString, "VSolveLinksKernel", sizeof(VSolveLinksCB) ); + vSolveLinksKernel = dxFunctions.compileComputeShaderFromString( VSolveLinksHLSLString, "VSolveLinksKernel", sizeof(VSolveLinksCB) ); if( !vSolveLinksKernel.constBuffer ) returnVal = false; - updateVelocitiesFromPositionsWithVelocitiesKernel = compileComputeShaderFromString( UpdateNodesHLSLString, "updateVelocitiesFromPositionsWithVelocitiesKernel", sizeof(UpdateVelocitiesFromPositionsWithVelocitiesCB) ); + updateVelocitiesFromPositionsWithVelocitiesKernel = dxFunctions.compileComputeShaderFromString( UpdateNodesHLSLString, "updateVelocitiesFromPositionsWithVelocitiesKernel", sizeof(UpdateVelocitiesFromPositionsWithVelocitiesCB) ); if( !updateVelocitiesFromPositionsWithVelocitiesKernel.constBuffer ) returnVal = false; - updateVelocitiesFromPositionsWithoutVelocitiesKernel = compileComputeShaderFromString( UpdatePositionsHLSLString, "updateVelocitiesFromPositionsWithoutVelocitiesKernel", sizeof(UpdateVelocitiesFromPositionsWithoutVelocitiesCB) ); + updateVelocitiesFromPositionsWithoutVelocitiesKernel = dxFunctions.compileComputeShaderFromString( UpdatePositionsHLSLString, "updateVelocitiesFromPositionsWithoutVelocitiesKernel", sizeof(UpdateVelocitiesFromPositionsWithoutVelocitiesCB) ); if( !updateVelocitiesFromPositionsWithoutVelocitiesKernel.constBuffer ) returnVal = false; - integrateKernel = compileComputeShaderFromString( IntegrateHLSLString, "IntegrateKernel", sizeof(IntegrateCB) ); + integrateKernel = dxFunctions.compileComputeShaderFromString( IntegrateHLSLString, "IntegrateKernel", sizeof(IntegrateCB) ); if( !integrateKernel.constBuffer ) returnVal = false; - applyForcesKernel = compileComputeShaderFromString( ApplyForcesHLSLString, "ApplyForcesKernel", sizeof(ApplyForcesCB) ); + applyForcesKernel = dxFunctions.compileComputeShaderFromString( ApplyForcesHLSLString, "ApplyForcesKernel", sizeof(ApplyForcesCB) ); if( !applyForcesKernel.constBuffer ) returnVal = false; + solveCollisionsAndUpdateVelocitiesKernel = dxFunctions.compileComputeShaderFromString( SolveCollisionsAndUpdateVelocitiesHLSLString, "SolveCollisionsAndUpdateVelocitiesKernel", sizeof(SolveCollisionsAndUpdateVelocitiesCB) ); + if( !solveCollisionsAndUpdateVelocitiesKernel.constBuffer ) + returnVal = false; // TODO: Rename to UpdateSoftBodies - resetNormalsAndAreasKernel = compileComputeShaderFromString( UpdateNormalsHLSLString, "ResetNormalsAndAreasKernel", sizeof(UpdateSoftBodiesCB) ); + resetNormalsAndAreasKernel = dxFunctions.compileComputeShaderFromString( UpdateNormalsHLSLString, "ResetNormalsAndAreasKernel", sizeof(UpdateSoftBodiesCB) ); if( !resetNormalsAndAreasKernel.constBuffer ) returnVal = false; - normalizeNormalsAndAreasKernel = compileComputeShaderFromString( UpdateNormalsHLSLString, "NormalizeNormalsAndAreasKernel", sizeof(UpdateSoftBodiesCB) ); + normalizeNormalsAndAreasKernel = dxFunctions.compileComputeShaderFromString( UpdateNormalsHLSLString, "NormalizeNormalsAndAreasKernel", sizeof(UpdateSoftBodiesCB) ); if( !normalizeNormalsAndAreasKernel.constBuffer ) returnVal = false; - updateSoftBodiesKernel = compileComputeShaderFromString( UpdateNormalsHLSLString, "UpdateSoftBodiesKernel", sizeof(UpdateSoftBodiesCB) ); + updateSoftBodiesKernel = dxFunctions.compileComputeShaderFromString( UpdateNormalsHLSLString, "UpdateSoftBodiesKernel", sizeof(UpdateSoftBodiesCB) ); if( !updateSoftBodiesKernel.constBuffer ) returnVal = false; - outputToVertexArrayWithNormalsKernel = compileComputeShaderFromString( OutputToVertexArrayHLSLString, "OutputToVertexArrayWithNormalsKernel", sizeof(OutputToVertexArrayCB) ); - if( !outputToVertexArrayWithNormalsKernel.constBuffer ) - returnVal = false; - outputToVertexArrayWithoutNormalsKernel = compileComputeShaderFromString( OutputToVertexArrayHLSLString, "OutputToVertexArrayWithoutNormalsKernel", sizeof(OutputToVertexArrayCB) ); - if( !outputToVertexArrayWithoutNormalsKernel.constBuffer ) + + computeBoundsKernel = dxFunctions.compileComputeShaderFromString( ComputeBoundsHLSLString, "ComputeBoundsKernel", sizeof(ComputeBoundsCB) ); + if( !computeBoundsKernel.constBuffer ) returnVal = false; @@ -1757,8 +2112,76 @@ bool btDX11SoftBodySolver::buildShaders() } +static Vectormath::Aos::Transform3 toTransform3( const btTransform &transform ) +{ + Vectormath::Aos::Transform3 outTransform; + outTransform.setCol(0, toVector3(transform.getBasis().getColumn(0))); + outTransform.setCol(1, toVector3(transform.getBasis().getColumn(1))); + outTransform.setCol(2, toVector3(transform.getBasis().getColumn(2))); + outTransform.setCol(3, toVector3(transform.getOrigin())); + return outTransform; +} + + +void btDX11SoftBodySolver::btAcceleratedSoftBodyInterface::updateBounds( const btVector3 &lowerBound, const btVector3 &upperBound ) +{ + float scalarMargin = this->getSoftBody()->getCollisionShape()->getMargin(); + btVector3 vectorMargin( scalarMargin, scalarMargin, scalarMargin ); + m_softBody->m_bounds[0] = lowerBound - vectorMargin; + m_softBody->m_bounds[1] = upperBound + vectorMargin; +} + +void btDX11SoftBodySolver::processCollision( btSoftBody*, btSoftBody* ) +{ + +} + +// Add the collision object to the set to deal with for a particular soft body +void btDX11SoftBodySolver::processCollision( btSoftBody *softBody, btCollisionObject* collisionObject ) +{ + int softBodyIndex = findSoftBodyIndex( softBody ); + + if( softBodyIndex >= 0 ) + { + btCollisionShape *collisionShape = collisionObject->getCollisionShape(); + float friction = collisionObject->getFriction(); + int shapeType = collisionShape->getShapeType(); + if( shapeType == CAPSULE_SHAPE_PROXYTYPE ) + { + // Add to the list of expected collision objects + CollisionShapeDescription newCollisionShapeDescription; + newCollisionShapeDescription.softBodyIdentifier = softBodyIndex; + newCollisionShapeDescription.collisionShapeType = shapeType; + // TODO: May need to transpose this matrix either here or in HLSL + newCollisionShapeDescription.shapeTransform = toTransform3(collisionObject->getWorldTransform()); + btCapsuleShape *capsule = static_cast( collisionShape ); + newCollisionShapeDescription.radius = capsule->getRadius(); + newCollisionShapeDescription.halfHeight = capsule->getHalfHeight(); + newCollisionShapeDescription.margin = capsule->getMargin(); + newCollisionShapeDescription.friction = friction; + btRigidBody* body = static_cast< btRigidBody* >( collisionObject ); + newCollisionShapeDescription.linearVelocity = toVector3(body->getLinearVelocity()); + newCollisionShapeDescription.angularVelocity = toVector3(body->getAngularVelocity()); + m_collisionObjectDetails.push_back( newCollisionShapeDescription ); + + } else { + btAssert("Unsupported collision shape type\n"); + } + } else { + btAssert("Unknown soft body"); + } +} // btDX11SoftBodySolver::processCollision + + + void btDX11SoftBodySolver::predictMotion( float timeStep ) { + // Clear the collision shape array for the next frame + // Ensure that the DX11 ones are moved off the device so they will be updated correctly + m_dx11CollisionObjectDetails.changedOnCPU(); + m_dx11PerClothCollisionObjects.changedOnCPU(); + m_collisionObjectDetails.clear(); + // Fill the force arrays with current acceleration data etc m_perClothWindVelocity.resize( m_softBodySet.size() ); for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) @@ -1774,6 +2197,12 @@ void btDX11SoftBodySolver::predictMotion( float timeStep ) // Itegrate motion for all soft bodies dealt with by the solver integrate( timeStep * getTimeScale() ); + + // Update bounds + // Will update the bounds for all softBodies being dealt with by the solver and + // set the values in the btSoftBody object + updateBounds(); + // End prediction work for solvers } diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11.h b/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11.h index ea5b3d462..f34d9f2cd 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11.h +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11.h @@ -26,184 +26,20 @@ subject to the following restrictions: -/** - * SoftBody class to maintain information about a soft body instance - * within a solver. - * This data addresses the main solver arrays. - */ -class btDX11AcceleratedSoftBodyInterface +class DXFunctions { -protected: - /** Current number of vertices that are part of this cloth */ - int m_numVertices; - /** Maximum number of vertices allocated to be part of this cloth */ - int m_maxVertices; - /** Current number of triangles that are part of this cloth */ - int m_numTriangles; - /** Maximum number of triangles allocated to be part of this cloth */ - int m_maxTriangles; - /** Index of first vertex in the world allocated to this cloth */ - int m_firstVertex; - /** Index of first triangle in the world allocated to this cloth */ - int m_firstTriangle; - /** Index of first link in the world allocated to this cloth */ - int m_firstLink; - /** Maximum number of links allocated to this cloth */ - int m_maxLinks; - /** Current number of links allocated to this cloth */ - int m_numLinks; - - /** The actual soft body this data represents */ - btSoftBody *m_softBody; - - public: - btDX11AcceleratedSoftBodyInterface( btSoftBody *softBody ) : - m_softBody( softBody ) - { - m_numVertices = 0; - m_maxVertices = 0; - m_numTriangles = 0; - m_maxTriangles = 0; - m_firstVertex = 0; - m_firstTriangle = 0; - m_firstLink = 0; - m_maxLinks = 0; - m_numLinks = 0; - } - int getNumVertices() - { - return m_numVertices; - } - - int getNumTriangles() - { - return m_numTriangles; - } - - int getMaxVertices() - { - return m_maxVertices; - } - - int getMaxTriangles() - { - return m_maxTriangles; - } - - int getFirstVertex() - { - return m_firstVertex; - } - - int getFirstTriangle() - { - return m_firstTriangle; - } - - // TODO: All of these set functions will have to do checks and - // update the world because restructuring of the arrays will be necessary - // Reasonable use of "friend"? - void setNumVertices( int numVertices ) - { - m_numVertices = numVertices; - } - void setNumTriangles( int numTriangles ) + ID3D11Device * m_dx11Device; + ID3D11DeviceContext* m_dx11Context; + + DXFunctions( ID3D11Device *dx11Device, ID3D11DeviceContext* dx11Context) : + m_dx11Device( dx11Device ), + m_dx11Context( dx11Context ) { - m_numTriangles = numTriangles; + } - void setMaxVertices( int maxVertices ) - { - m_maxVertices = maxVertices; - } - - void setMaxTriangles( int maxTriangles ) - { - m_maxTriangles = maxTriangles; - } - - void setFirstVertex( int firstVertex ) - { - m_firstVertex = firstVertex; - } - - void setFirstTriangle( int firstTriangle ) - { - m_firstTriangle = firstTriangle; - } - - void setMaxLinks( int maxLinks ) - { - m_maxLinks = maxLinks; - } - - void setNumLinks( int numLinks ) - { - m_numLinks = numLinks; - } - - void setFirstLink( int firstLink ) - { - m_firstLink = firstLink; - } - - int getMaxLinks() - { - return m_maxLinks; - } - - int getNumLinks() - { - return m_numLinks; - } - - int getFirstLink() - { - return m_firstLink; - } - - btSoftBody* getSoftBody() - { - return m_softBody; - } - -#if 0 - void setAcceleration( Vectormath::Aos::Vector3 acceleration ) - { - m_currentSolver->setPerClothAcceleration( m_clothIdentifier, acceleration ); - } - - void setWindVelocity( Vectormath::Aos::Vector3 windVelocity ) - { - m_currentSolver->setPerClothWindVelocity( m_clothIdentifier, windVelocity ); - } - - /** - * Set the density of the air in which the cloth is situated. - */ - void setAirDensity( btScalar density ) - { - m_currentSolver->setPerClothMediumDensity( m_clothIdentifier, static_cast(density) ); - } - - /** - * Add a collision object to this soft body. - */ - void addCollisionObject( btCollisionObject *collisionObject ) - { - m_currentSolver->addCollisionObjectForSoftBody( m_clothIdentifier, collisionObject ); - } -#endif -}; - - -class btDX11SoftBodySolver : public btSoftBodySolver -{ -public: - - class KernelDesc { protected: @@ -226,6 +62,238 @@ public: } }; + /** + * Compile a compute shader kernel from a string and return the appropriate KernelDesc object. + */ + KernelDesc compileComputeShaderFromString( const char* shaderString, const char* shaderName, int constBufferSize, D3D10_SHADER_MACRO *compileMacros = 0 ); + +}; + +class btDX11SoftBodySolver : public btSoftBodySolver +{ +protected: + /** + * Entry in the collision shape array. + * Specifies the shape type, the transform matrix and the necessary details of the collisionShape. + */ + struct CollisionShapeDescription + { + Vectormath::Aos::Transform3 shapeTransform; + Vectormath::Aos::Vector3 linearVelocity; + Vectormath::Aos::Vector3 angularVelocity; + + int softBodyIdentifier; + int collisionShapeType; + + // Both needed for capsule + float radius; + float halfHeight; + + float margin; + float friction; + + CollisionShapeDescription() + { + collisionShapeType = 0; + margin = 0; + friction = 0; + } + }; + + struct UIntVector3 + { + UIntVector3() + { + x = 0; + y = 0; + z = 0; + _padding = 0; + } + + UIntVector3( unsigned int x_, unsigned int y_, unsigned int z_ ) + { + x = x_; + y = y_; + z = z_; + _padding = 0; + } + + unsigned int x; + unsigned int y; + unsigned int z; + unsigned int _padding; + }; + + + +public: + /** + * SoftBody class to maintain information about a soft body instance + * within a solver. + * This data addresses the main solver arrays. + */ + class btAcceleratedSoftBodyInterface + { + protected: + /** Current number of vertices that are part of this cloth */ + int m_numVertices; + /** Maximum number of vertices allocated to be part of this cloth */ + int m_maxVertices; + /** Current number of triangles that are part of this cloth */ + int m_numTriangles; + /** Maximum number of triangles allocated to be part of this cloth */ + int m_maxTriangles; + /** Index of first vertex in the world allocated to this cloth */ + int m_firstVertex; + /** Index of first triangle in the world allocated to this cloth */ + int m_firstTriangle; + /** Index of first link in the world allocated to this cloth */ + int m_firstLink; + /** Maximum number of links allocated to this cloth */ + int m_maxLinks; + /** Current number of links allocated to this cloth */ + int m_numLinks; + + /** The actual soft body this data represents */ + btSoftBody *m_softBody; + + + public: + btAcceleratedSoftBodyInterface( btSoftBody *softBody ) : + m_softBody( softBody ) + { + m_numVertices = 0; + m_maxVertices = 0; + m_numTriangles = 0; + m_maxTriangles = 0; + m_firstVertex = 0; + m_firstTriangle = 0; + m_firstLink = 0; + m_maxLinks = 0; + m_numLinks = 0; + } + int getNumVertices() const + { + return m_numVertices; + } + + int getNumTriangles() const + { + return m_numTriangles; + } + + int getMaxVertices() const + { + return m_maxVertices; + } + + int getMaxTriangles() const + { + return m_maxTriangles; + } + + int getFirstVertex() const + { + return m_firstVertex; + } + + int getFirstTriangle() const + { + return m_firstTriangle; + } + + + /** + * Update the bounds in the btSoftBody object + */ + void updateBounds( const btVector3 &lowerBound, const btVector3 &upperBound ); + + // TODO: All of these set functions will have to do checks and + // update the world because restructuring of the arrays will be necessary + // Reasonable use of "friend"? + void setNumVertices( int numVertices ) + { + m_numVertices = numVertices; + } + + void setNumTriangles( int numTriangles ) + { + m_numTriangles = numTriangles; + } + + void setMaxVertices( int maxVertices ) + { + m_maxVertices = maxVertices; + } + + void setMaxTriangles( int maxTriangles ) + { + m_maxTriangles = maxTriangles; + } + + void setFirstVertex( int firstVertex ) + { + m_firstVertex = firstVertex; + } + + void setFirstTriangle( int firstTriangle ) + { + m_firstTriangle = firstTriangle; + } + + void setMaxLinks( int maxLinks ) + { + m_maxLinks = maxLinks; + } + + void setNumLinks( int numLinks ) + { + m_numLinks = numLinks; + } + + void setFirstLink( int firstLink ) + { + m_firstLink = firstLink; + } + + int getMaxLinks() + { + return m_maxLinks; + } + + int getNumLinks() + { + return m_numLinks; + } + + int getFirstLink() + { + return m_firstLink; + } + + btSoftBody* getSoftBody() + { + return m_softBody; + } + + }; + + + struct CollisionObjectIndices + { + CollisionObjectIndices( int f, int e ) + { + firstObject = f; + endObject = e; + } + + int firstObject; + int endObject; + }; + + + + struct PrepareLinksCB { @@ -284,20 +352,6 @@ public: }; - struct OutputToVertexArrayCB - { - int startNode; - int numNodes; - int positionOffset; - int positionStride; - - int normalOffset; - int normalStride; - int padding1; - int padding2; - }; - - struct ApplyForcesCB { unsigned int numNodes; @@ -326,17 +380,38 @@ public: int padding; }; + struct ComputeBoundsCB + { + int numNodes; + int numSoftBodies; + int padding1; + int padding2; + }; -private: + struct SolveCollisionsAndUpdateVelocitiesCB + { + unsigned int numNodes; + float isolverdt; + int padding0; + int padding1; + }; + + + + +protected: ID3D11Device * m_dx11Device; ID3D11DeviceContext* m_dx11Context; - - + + DXFunctions dxFunctions; +public: /** Link data for all cloths. Note that this will be sorted batch-wise for efficient computation and m_linkAddresses will maintain the addressing. */ btSoftBodyLinkDataDX11 m_linkData; btSoftBodyVertexDataDX11 m_vertexData; btSoftBodyTriangleDataDX11 m_triangleData; - + +protected: + /** Variable to define whether we need to update solver constants on the next iteration */ bool m_updateSolverConstants; @@ -346,7 +421,7 @@ private: * Cloths owned by this solver. * Only our cloths are in this array. */ - btAlignedObjectArray< btDX11AcceleratedSoftBodyInterface * > m_softBodySet; + btAlignedObjectArray< btAcceleratedSoftBodyInterface * > m_softBodySet; /** Acceleration value to be applied to all non-static vertices in the solver. * Index n is cloth n, array sized by number of cloths in the world not the solver. @@ -380,24 +455,59 @@ private: btAlignedObjectArray< float > m_perClothMediumDensity; btDX11Buffer m_dx11PerClothMediumDensity; - KernelDesc prepareLinksKernel; - KernelDesc solvePositionsFromLinksKernel; - KernelDesc vSolveLinksKernel; - KernelDesc integrateKernel; - KernelDesc addVelocityKernel; - KernelDesc updatePositionsFromVelocitiesKernel; - KernelDesc updateVelocitiesFromPositionsWithoutVelocitiesKernel; - KernelDesc updateVelocitiesFromPositionsWithVelocitiesKernel; - KernelDesc resetNormalsAndAreasKernel; - KernelDesc normalizeNormalsAndAreasKernel; - KernelDesc updateSoftBodiesKernel; - KernelDesc outputToVertexArrayWithNormalsKernel; - KernelDesc outputToVertexArrayWithoutNormalsKernel; + + /** + * Collision shape details: pair of index of first collision shape for the cloth and number of collision objects. + */ + btAlignedObjectArray< CollisionObjectIndices > m_perClothCollisionObjects; + btDX11Buffer m_dx11PerClothCollisionObjects; - KernelDesc outputToVertexArrayKernel; - KernelDesc applyForcesKernel; - KernelDesc collideSphereKernel; - KernelDesc collideCylinderKernel; + /** + * Collision shapes being passed across to the cloths in this solver. + */ + btAlignedObjectArray< CollisionShapeDescription > m_collisionObjectDetails; + btDX11Buffer< CollisionShapeDescription > m_dx11CollisionObjectDetails; + + /** + * Minimum bounds for each cloth. + * Updated by GPU and returned for use by broad phase. + * These are int vectors as a reminder that they store the int representation of a float, not a float. + * Bit 31 is inverted - is floats are stored with int-sortable values. + */ + btAlignedObjectArray< UIntVector3 > m_perClothMinBounds; + btDX11Buffer< UIntVector3 > m_dx11PerClothMinBounds; + + /** + * Maximum bounds for each cloth. + * Updated by GPU and returned for use by broad phase. + * These are int vectors as a reminder that they store the int representation of a float, not a float. + * Bit 31 is inverted - is floats are stored with int-sortable values. + */ + btAlignedObjectArray< UIntVector3 > m_perClothMaxBounds; + btDX11Buffer< UIntVector3 > m_dx11PerClothMaxBounds; + + + /** + * Friction coefficient for each cloth + */ + btAlignedObjectArray< float > m_perClothFriction; + btDX11Buffer< float > m_dx11PerClothFriction; + + DXFunctions::KernelDesc prepareLinksKernel; + DXFunctions::KernelDesc solvePositionsFromLinksKernel; + DXFunctions::KernelDesc vSolveLinksKernel; + DXFunctions::KernelDesc integrateKernel; + DXFunctions::KernelDesc addVelocityKernel; + DXFunctions::KernelDesc updatePositionsFromVelocitiesKernel; + DXFunctions::KernelDesc updateVelocitiesFromPositionsWithoutVelocitiesKernel; + DXFunctions::KernelDesc updateVelocitiesFromPositionsWithVelocitiesKernel; + DXFunctions::KernelDesc solveCollisionsAndUpdateVelocitiesKernel; + DXFunctions::KernelDesc resetNormalsAndAreasKernel; + DXFunctions::KernelDesc normalizeNormalsAndAreasKernel; + DXFunctions::KernelDesc computeBoundsKernel; + DXFunctions::KernelDesc updateSoftBodiesKernel; + + DXFunctions::KernelDesc applyForcesKernel; /** @@ -410,12 +520,7 @@ private: const Vectormath::Aos::Point3 &vertex2 ); - /** - * Compile a compute shader kernel from a string and return the appropriate KernelDesc object. - */ - KernelDesc compileComputeShaderFromString( const char* shaderString, const char* shaderName, int constBufferSize ); - - bool buildShaders(); + virtual bool buildShaders(); void resetNormalsAndAreas( int numVertices ); @@ -423,19 +528,20 @@ private: void executeUpdateSoftBodies( int firstTriangle, int numTriangles ); + void prepareCollisionConstraints(); + Vectormath::Aos::Vector3 ProjectOnAxis( const Vectormath::Aos::Vector3 &v, const Vectormath::Aos::Vector3 &a ); void ApplyClampedForce( float solverdt, const Vectormath::Aos::Vector3 &force, const Vectormath::Aos::Vector3 &vertexVelocity, float inverseMass, Vectormath::Aos::Vector3 &vertexForce ); virtual void applyForces( float solverdt ); - void updateConstants( float timeStep ); - - btDX11AcceleratedSoftBodyInterface *findSoftBodyInterface( const btSoftBody* const softBody ); + virtual void updateConstants( float timeStep ); + int findSoftBodyIndex( const btSoftBody* const softBody ); ////////////////////////////////////// // Kernel dispatches - void prepareLinks(); + virtual void prepareLinks(); void updatePositionsFromVelocities( float solverdt ); void solveLinksForPosition( int startLink, int numLinks, float kst, float ti ); @@ -443,14 +549,27 @@ private: void updateVelocitiesFromPositionsWithVelocities( float isolverdt ); void updateVelocitiesFromPositionsWithoutVelocities( float isolverdt ); + void computeBounds( ); + void solveCollisionsAndUpdateVelocities( float isolverdt ); // End kernel dispatches ///////////////////////////////////// + void updateBounds(); + + + void releaseKernels(); + public: btDX11SoftBodySolver(ID3D11Device * dx11Device, ID3D11DeviceContext* dx11Context); virtual ~btDX11SoftBodySolver(); + + + virtual SolverTypes getSolverType() const + { + return DX_SOLVER; + } virtual btSoftBodyLinkData &getLinkData(); @@ -461,19 +580,94 @@ public: + + + btAcceleratedSoftBodyInterface *findSoftBodyInterface( const btSoftBody* const softBody ); + const btAcceleratedSoftBodyInterface * const findSoftBodyInterface( const btSoftBody* const softBody ) const; virtual bool checkInitialized(); virtual void updateSoftBodies( ); - virtual void optimize( btAlignedObjectArray< btSoftBody * > &softBodies ); + virtual void optimize( btAlignedObjectArray< btSoftBody * > &softBodies , bool forceUpdate=false); + + virtual void copyBackToSoftBodies(); virtual void solveConstraints( float solverdt ); virtual void predictMotion( float solverdt ); - virtual void copySoftBodyToVertexBuffer( const btSoftBody *const softBody, btVertexBufferDescriptor *vertexBuffer ); + + virtual void processCollision( btSoftBody *, btCollisionObject* ); + virtual void processCollision( btSoftBody*, btSoftBody* ); + +}; + + + +/** + * Class to manage movement of data from a solver to a given target. + * This version is the DX to CPU version. + */ +class btSoftBodySolverOutputDXtoCPU : public btSoftBodySolverOutput +{ +protected: + +public: + btSoftBodySolverOutputDXtoCPU() + { + } + + /** Output current computed vertex data to the vertex buffers for all cloths in the solver. */ + virtual void copySoftBodyToVertexBuffer( const btSoftBody * const softBody, btVertexBufferDescriptor *vertexBuffer ); +}; + +/** + * Class to manage movement of data from a solver to a given target. + * This version is the DX to DX version and subclasses DX to CPU so that it works for that too. + */ +class btSoftBodySolverOutputDXtoDX : public btSoftBodySolverOutputDXtoCPU +{ +protected: + struct OutputToVertexArrayCB + { + int startNode; + int numNodes; + int positionOffset; + int positionStride; + + int normalOffset; + int normalStride; + int padding1; + int padding2; + }; + + DXFunctions dxFunctions; + DXFunctions::KernelDesc outputToVertexArrayWithNormalsKernel; + DXFunctions::KernelDesc outputToVertexArrayWithoutNormalsKernel; + + + bool m_shadersInitialized; + + bool checkInitialized(); + bool buildShaders(); + void releaseKernels(); + +public: + btSoftBodySolverOutputDXtoDX(ID3D11Device *dx11Device, ID3D11DeviceContext* dx11Context) : + dxFunctions( dx11Device, dx11Context ) + { + m_shadersInitialized = false; + } + + ~btSoftBodySolverOutputDXtoDX() + { + releaseKernels(); + } + + /** Output current computed vertex data to the vertex buffers for all cloths in the solver. */ + virtual void copySoftBodyToVertexBuffer( const btSoftBody * const softBody, btVertexBufferDescriptor *vertexBuffer ); }; #endif // #ifndef BT_ACCELERATED_SOFT_BODY_DX11_SOLVER_H diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11SIMDAware.cpp b/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11SIMDAware.cpp index c72dead3e..896529d95 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11SIMDAware.cpp +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11SIMDAware.cpp @@ -18,6 +18,7 @@ subject to the following restrictions: #define WAVEFRONT_SIZE 32 #define WAVEFRONT_BLOCK_MULTIPLIER 2 +#define GROUP_SIZE (WAVEFRONT_SIZE*WAVEFRONT_BLOCK_MULTIPLIER) #define LINKS_PER_SIMD_LANE 16 #define STRINGIFY( S ) STRINGIFY2( S ) @@ -30,10 +31,9 @@ subject to the following restrictions: #include "btSoftBodySolver_DX11SIMDAware.h" #include "btSoftBodySolverVertexBuffer_DX11.h" #include "BulletSoftBody/btSoftBody.h" +#include "BulletCollision/CollisionShapes/btCapsuleShape.h" #define MSTRINGIFY(A) #A -static char* PrepareLinksHLSLString = -#include "HLSL/PrepareLinks.hlsl" static char* UpdatePositionsFromVelocitiesHLSLString = #include "HLSL/UpdatePositionsFromVelocities.hlsl" static char* SolvePositionsSIMDBatchedHLSLString = @@ -54,6 +54,10 @@ static char* OutputToVertexArrayHLSLString = #include "HLSL/OutputToVertexArray.hlsl" static char* VSolveLinksHLSLString = #include "HLSL/VSolveLinks.hlsl" +static char* ComputeBoundsHLSLString = +#include "HLSL/ComputeBounds.hlsl" +static char* SolveCollisionsAndUpdateVelocitiesHLSLString = +#include "HLSL/solveCollisionsAndUpdateVelocitiesSIMDBatched.hlsl" @@ -170,18 +174,8 @@ bool btSoftBodyLinkDataDX11SIMDAware::moveFromAccelerator() btDX11SIMDAwareSoftBodySolver::btDX11SIMDAwareSoftBodySolver(ID3D11Device * dx11Device, ID3D11DeviceContext* dx11Context) : - m_dx11Device( dx11Device ), - m_dx11Context( dx11Context ), - m_linkData(m_dx11Device, m_dx11Context), - m_vertexData(m_dx11Device, m_dx11Context), - m_triangleData(m_dx11Device, m_dx11Context), - m_dx11PerClothAcceleration( m_dx11Device, m_dx11Context, &m_perClothAcceleration, true ), - m_dx11PerClothWindVelocity( m_dx11Device, m_dx11Context, &m_perClothWindVelocity, true ), - m_dx11PerClothDampingFactor( m_dx11Device, m_dx11Context, &m_perClothDampingFactor, true ), - m_dx11PerClothVelocityCorrectionCoefficient( m_dx11Device, m_dx11Context, &m_perClothVelocityCorrectionCoefficient, true ), - m_dx11PerClothLiftFactor( m_dx11Device, m_dx11Context, &m_perClothLiftFactor, true ), - m_dx11PerClothDragFactor( m_dx11Device, m_dx11Context, &m_perClothDragFactor, true ), - m_dx11PerClothMediumDensity( m_dx11Device, m_dx11Context, &m_perClothMediumDensity, true ) + btDX11SoftBodySolver( dx11Device, dx11Context ), + m_linkData(m_dx11Device, m_dx11Context) { // Initial we will clearly need to update solver constants // For now this is global for the cloths linked with this solver - we should probably make this body specific @@ -191,48 +185,20 @@ btDX11SIMDAwareSoftBodySolver::btDX11SIMDAwareSoftBodySolver(ID3D11Device * dx11 m_shadersInitialized = false; } -void btDX11SIMDAwareSoftBodySolver::releaseKernels() -{ - SAFE_RELEASE( integrateKernel.constBuffer ); - SAFE_RELEASE( integrateKernel.kernel ); - SAFE_RELEASE( solvePositionsFromLinksKernel.constBuffer ); - SAFE_RELEASE( solvePositionsFromLinksKernel.kernel ); - SAFE_RELEASE( updatePositionsFromVelocitiesKernel.constBuffer ); - SAFE_RELEASE( updatePositionsFromVelocitiesKernel.kernel ); - SAFE_RELEASE( updateVelocitiesFromPositionsWithoutVelocitiesKernel.constBuffer ); - SAFE_RELEASE( updateVelocitiesFromPositionsWithoutVelocitiesKernel.kernel ); - SAFE_RELEASE( updateVelocitiesFromPositionsWithVelocitiesKernel.constBuffer ); - SAFE_RELEASE( updateVelocitiesFromPositionsWithVelocitiesKernel.kernel ); - SAFE_RELEASE( resetNormalsAndAreasKernel.constBuffer ); - SAFE_RELEASE( resetNormalsAndAreasKernel.kernel ); - SAFE_RELEASE( normalizeNormalsAndAreasKernel.constBuffer ); - SAFE_RELEASE( normalizeNormalsAndAreasKernel.kernel ); - SAFE_RELEASE( updateSoftBodiesKernel.constBuffer ); - SAFE_RELEASE( updateSoftBodiesKernel.kernel ); - SAFE_RELEASE( outputToVertexArrayWithNormalsKernel.constBuffer ); - SAFE_RELEASE( outputToVertexArrayWithNormalsKernel.kernel ); - SAFE_RELEASE( outputToVertexArrayWithoutNormalsKernel.constBuffer ); - SAFE_RELEASE( outputToVertexArrayWithoutNormalsKernel.kernel ); - - - SAFE_RELEASE( addVelocityKernel.constBuffer ); - SAFE_RELEASE( addVelocityKernel.kernel ); - SAFE_RELEASE( applyForcesKernel.constBuffer ); - SAFE_RELEASE( applyForcesKernel.kernel ); - SAFE_RELEASE( outputToVertexArrayKernel.constBuffer ); - SAFE_RELEASE( outputToVertexArrayKernel.kernel ); - SAFE_RELEASE( collideCylinderKernel.constBuffer ); - SAFE_RELEASE( collideCylinderKernel.kernel ); - - m_shadersInitialized = false; -} - btDX11SIMDAwareSoftBodySolver::~btDX11SIMDAwareSoftBodySolver() { releaseKernels(); } +btSoftBodyLinkData &btDX11SIMDAwareSoftBodySolver::getLinkData() +{ + // TODO: Consider setting link data to "changed" here + return m_linkData; +} + + + void btDX11SIMDAwareSoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * > &softBodies ) { if( m_softBodySet.size() != softBodies.size() ) @@ -260,14 +226,21 @@ void btDX11SIMDAwareSoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * m_perClothLiftFactor.push_back( softBody->m_cfg.kLF ); m_perClothDragFactor.push_back( softBody->m_cfg.kDG ); m_perClothMediumDensity.push_back(softBody->getWorldInfo()->air_density); + // Simple init values. Actually we'll put 0 and -1 into them at the appropriate time + m_perClothMinBounds.push_back( UIntVector3( 0, 0, 0 ) ); + m_perClothMaxBounds.push_back( UIntVector3( UINT_MAX, UINT_MAX, UINT_MAX ) ); + m_perClothFriction.push_back( softBody->getFriction() ); + m_perClothCollisionObjects.push_back( CollisionObjectIndices(-1, -1) ); // Add space for new vertices and triangles in the default solver for now // TODO: Include space here for tearing too later int firstVertex = getVertexData().getNumVertices(); int numVertices = softBody->m_nodes.size(); - int maxVertices = numVertices; + // Round maxVertices to a multiple of the workgroup size so we know we're safe to run over in a given group + // maxVertices can be increased to allow tearing, but should be used sparingly because these extra verts will always be processed + int maxVertices = GROUP_SIZE*((numVertices+GROUP_SIZE)/GROUP_SIZE); // Allocate space for new vertices in all the vertex arrays - getVertexData().createVertices( maxVertices, softBodyIndex ); + getVertexData().createVertices( numVertices, softBodyIndex, maxVertices ); int firstTriangle = getTriangleData().getNumTriangles(); int numTriangles = softBody->m_faces.size(); @@ -352,341 +325,61 @@ void btDX11SIMDAwareSoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * } -btSoftBodyLinkData &btDX11SIMDAwareSoftBodySolver::getLinkData() + +void btDX11SIMDAwareSoftBodySolver::solveConstraints( float solverdt ) { - // TODO: Consider setting link data to "changed" here - return m_linkData; -} - -btSoftBodyVertexData &btDX11SIMDAwareSoftBodySolver::getVertexData() -{ - // TODO: Consider setting vertex data to "changed" here - return m_vertexData; -} - -btSoftBodyTriangleData &btDX11SIMDAwareSoftBodySolver::getTriangleData() -{ - // TODO: Consider setting triangle data to "changed" here - return m_triangleData; -} - - -bool btDX11SIMDAwareSoftBodySolver::checkInitialized() -{ - if( !m_shadersInitialized ) - if( buildShaders() ) - m_shadersInitialized = true; - - return m_shadersInitialized; -} - -void btDX11SIMDAwareSoftBodySolver::resetNormalsAndAreas( int numVertices ) -{ - // No need to batch link solver, it is entirely parallel - // Copy kernel parameters to GPU - UpdateSoftBodiesCB constBuffer; - - constBuffer.numNodes = numVertices; - constBuffer.epsilon = FLT_EPSILON; - - // Todo: factor this out. Number of nodes is static and sdt might be, too, we can update this just once on setup - D3D11_MAPPED_SUBRESOURCE MappedResource = {0}; - m_dx11Context->Map( integrateKernel.constBuffer, 0, D3D11_MAP_WRITE_DISCARD, 0, &MappedResource ); - memcpy( MappedResource.pData, &constBuffer, sizeof(UpdateSoftBodiesCB) ); - m_dx11Context->Unmap( integrateKernel.constBuffer, 0 ); - m_dx11Context->CSSetConstantBuffers( 0, 1, &integrateKernel.constBuffer ); - - // Set resources and dispatch - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &(m_vertexData.m_dx11VertexNormal.getUAV()), NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &(m_vertexData.m_dx11VertexArea.getUAV()), NULL ); - - // Execute the kernel - m_dx11Context->CSSetShader( resetNormalsAndAreasKernel.kernel, NULL, 0 ); - - int numBlocks = (constBuffer.numNodes + (128-1)) / 128; - m_dx11Context->Dispatch(numBlocks, 1, 1 ); - - { - // Tidy up - ID3D11UnorderedAccessView* pUAViewNULL = NULL; - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &pUAViewNULL, NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &pUAViewNULL, NULL ); - - ID3D11Buffer *pBufferNull = NULL; - m_dx11Context->CSSetConstantBuffers( 0, 1, &pBufferNull ); - } -} // btDX11SIMDAwareSoftBodySolver::resetNormalsAndAreas - -void btDX11SIMDAwareSoftBodySolver::normalizeNormalsAndAreas( int numVertices ) -{ - // No need to batch link solver, it is entirely parallel - // Copy kernel parameters to GPU - UpdateSoftBodiesCB constBuffer; - - constBuffer.numNodes = numVertices; - constBuffer.epsilon = FLT_EPSILON; - - // Todo: factor this out. Number of nodes is static and sdt might be, too, we can update this just once on setup - D3D11_MAPPED_SUBRESOURCE MappedResource = {0}; - m_dx11Context->Map( integrateKernel.constBuffer, 0, D3D11_MAP_WRITE_DISCARD, 0, &MappedResource ); - memcpy( MappedResource.pData, &constBuffer, sizeof(UpdateSoftBodiesCB) ); - m_dx11Context->Unmap( integrateKernel.constBuffer, 0 ); - m_dx11Context->CSSetConstantBuffers( 0, 1, &integrateKernel.constBuffer ); - - // Set resources and dispatch - m_dx11Context->CSSetShaderResources( 2, 1, &(m_vertexData.m_dx11VertexTriangleCount.getSRV()) ); - - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &(m_vertexData.m_dx11VertexNormal.getUAV()), NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &(m_vertexData.m_dx11VertexArea.getUAV()), NULL ); - - // Execute the kernel - m_dx11Context->CSSetShader( normalizeNormalsAndAreasKernel.kernel, NULL, 0 ); - - int numBlocks = (constBuffer.numNodes + (128-1)) / 128; - m_dx11Context->Dispatch(numBlocks, 1, 1 ); - - { - // Tidy up - ID3D11ShaderResourceView* pViewNULL = NULL; - m_dx11Context->CSSetShaderResources( 2, 1, &pViewNULL ); - - ID3D11UnorderedAccessView* pUAViewNULL = NULL; - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &pUAViewNULL, NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &pUAViewNULL, NULL ); - - ID3D11Buffer *pBufferNull = NULL; - m_dx11Context->CSSetConstantBuffers( 0, 1, &pBufferNull ); - } -} // btDX11SIMDAwareSoftBodySolver::normalizeNormalsAndAreas - -void btDX11SIMDAwareSoftBodySolver::executeUpdateSoftBodies( int firstTriangle, int numTriangles ) -{ - // No need to batch link solver, it is entirely parallel - // Copy kernel parameters to GPU - UpdateSoftBodiesCB constBuffer; - - constBuffer.startFace = firstTriangle; - constBuffer.numFaces = numTriangles; - - // Todo: factor this out. Number of nodes is static and sdt might be, too, we can update this just once on setup - D3D11_MAPPED_SUBRESOURCE MappedResource = {0}; - m_dx11Context->Map( updateSoftBodiesKernel.constBuffer, 0, D3D11_MAP_WRITE_DISCARD, 0, &MappedResource ); - memcpy( MappedResource.pData, &constBuffer, sizeof(UpdateSoftBodiesCB) ); - m_dx11Context->Unmap( updateSoftBodiesKernel.constBuffer, 0 ); - m_dx11Context->CSSetConstantBuffers( 0, 1, &updateSoftBodiesKernel.constBuffer ); - - // Set resources and dispatch - m_dx11Context->CSSetShaderResources( 0, 1, &(m_triangleData.m_dx11VertexIndices.getSRV()) ); - m_dx11Context->CSSetShaderResources( 1, 1, &(m_vertexData.m_dx11VertexPosition.getSRV()) ); - - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &(m_vertexData.m_dx11VertexNormal.getUAV()), NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &(m_vertexData.m_dx11VertexArea.getUAV()), NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 2, 1, &(m_triangleData.m_dx11Normal.getUAV()), NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 3, 1, &(m_triangleData.m_dx11Area.getUAV()), NULL ); - - // Execute the kernel - m_dx11Context->CSSetShader( updateSoftBodiesKernel.kernel, NULL, 0 ); - - int numBlocks = (numTriangles + (128-1)) / 128; - m_dx11Context->Dispatch(numBlocks, 1, 1 ); - - { - // Tidy up - ID3D11ShaderResourceView* pViewNULL = NULL; - m_dx11Context->CSSetShaderResources( 4, 1, &pViewNULL ); - - ID3D11UnorderedAccessView* pUAViewNULL = NULL; - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &pUAViewNULL, NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &pUAViewNULL, NULL ); - - ID3D11Buffer *pBufferNull = NULL; - m_dx11Context->CSSetConstantBuffers( 0, 1, &pBufferNull ); - } -} // btDX11SIMDAwareSoftBodySolver::executeUpdateSoftBodies - -void btDX11SIMDAwareSoftBodySolver::updateSoftBodies() -{ - using namespace Vectormath::Aos; + //std::cerr << "'GPU' solve constraints\n"; + using Vectormath::Aos::Vector3; + using Vectormath::Aos::Point3; + using Vectormath::Aos::lengthSqr; + using Vectormath::Aos::dot; + // Prepare links + int numLinks = m_linkData.getNumLinks(); int numVertices = m_vertexData.getNumVertices(); - int numTriangles = m_triangleData.getNumTriangles(); + + float kst = 1.f; + float ti = 0.f; + + + m_dx11PerClothDampingFactor.moveToGPU(); + m_dx11PerClothVelocityCorrectionCoefficient.moveToGPU(); + + // Ensure data is on accelerator - m_vertexData.moveToAccelerator(); - m_triangleData.moveToAccelerator(); - - resetNormalsAndAreas( numVertices ); - - - // Go through triangle batches so updates occur correctly - for( int batchIndex = 0; batchIndex < m_triangleData.m_batchStartLengths.size(); ++batchIndex ) - { - - int startTriangle = m_triangleData.m_batchStartLengths[batchIndex].start; - int numTriangles = m_triangleData.m_batchStartLengths[batchIndex].length; - - executeUpdateSoftBodies( startTriangle, numTriangles ); - } - - - normalizeNormalsAndAreas( numVertices ); - -} // btDX11SIMDAwareSoftBodySolver::updateSoftBodies - - -Vectormath::Aos::Vector3 btDX11SIMDAwareSoftBodySolver::ProjectOnAxis( const Vectormath::Aos::Vector3 &v, const Vectormath::Aos::Vector3 &a ) -{ - return a*Vectormath::Aos::dot(v, a); -} - -void btDX11SIMDAwareSoftBodySolver::ApplyClampedForce( float solverdt, const Vectormath::Aos::Vector3 &force, const Vectormath::Aos::Vector3 &vertexVelocity, float inverseMass, Vectormath::Aos::Vector3 &vertexForce ) -{ - float dtInverseMass = solverdt*inverseMass; - if( Vectormath::Aos::lengthSqr(force * dtInverseMass) > Vectormath::Aos::lengthSqr(vertexVelocity) ) - { - vertexForce -= ProjectOnAxis( vertexVelocity, normalize( force ) )/dtInverseMass; - } else { - vertexForce += force; - } -} - -void btDX11SIMDAwareSoftBodySolver::applyForces( float solverdt ) -{ - using namespace Vectormath::Aos; - - // Ensure data is on accelerator - m_vertexData.moveToAccelerator(); - m_dx11PerClothAcceleration.moveToGPU(); - m_dx11PerClothLiftFactor.moveToGPU(); - m_dx11PerClothDragFactor.moveToGPU(); - m_dx11PerClothMediumDensity.moveToGPU(); - m_dx11PerClothWindVelocity.moveToGPU(); - - // No need to batch link solver, it is entirely parallel - // Copy kernel parameters to GPU - ApplyForcesCB constBuffer; - - constBuffer.numNodes = m_vertexData.getNumVertices(); - constBuffer.solverdt = solverdt; - constBuffer.epsilon = FLT_EPSILON; - - // Todo: factor this out. Number of nodes is static and sdt might be, too, we can update this just once on setup - D3D11_MAPPED_SUBRESOURCE MappedResource = {0}; - m_dx11Context->Map( integrateKernel.constBuffer, 0, D3D11_MAP_WRITE_DISCARD, 0, &MappedResource ); - memcpy( MappedResource.pData, &constBuffer, sizeof(ApplyForcesCB) ); - m_dx11Context->Unmap( integrateKernel.constBuffer, 0 ); - m_dx11Context->CSSetConstantBuffers( 0, 1, &integrateKernel.constBuffer ); - - // Set resources and dispatch - m_dx11Context->CSSetShaderResources( 0, 1, &(m_vertexData.m_dx11ClothIdentifier.getSRV()) ); - m_dx11Context->CSSetShaderResources( 1, 1, &(m_vertexData.m_dx11VertexNormal.getSRV()) ); - m_dx11Context->CSSetShaderResources( 2, 1, &(m_vertexData.m_dx11VertexArea.getSRV()) ); - m_dx11Context->CSSetShaderResources( 3, 1, &(m_vertexData.m_dx11VertexInverseMass.getSRV()) ); - m_dx11Context->CSSetShaderResources( 4, 1, &(m_dx11PerClothLiftFactor.getSRV()) ); - m_dx11Context->CSSetShaderResources( 5, 1, &(m_dx11PerClothDragFactor.getSRV()) ); - m_dx11Context->CSSetShaderResources( 6, 1, &(m_dx11PerClothWindVelocity.getSRV()) ); - m_dx11Context->CSSetShaderResources( 7, 1, &(m_dx11PerClothAcceleration.getSRV()) ); - m_dx11Context->CSSetShaderResources( 8, 1, &(m_dx11PerClothMediumDensity.getSRV()) ); - - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &(m_vertexData.m_dx11VertexForceAccumulator.getUAV()), NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &(m_vertexData.m_dx11VertexVelocity.getUAV()), NULL ); - - // Execute the kernel - m_dx11Context->CSSetShader( applyForcesKernel.kernel, NULL, 0 ); - - int numBlocks = (constBuffer.numNodes + (128-1)) / 128; - m_dx11Context->Dispatch(numBlocks, 1, 1 ); - - { - // Tidy up - ID3D11ShaderResourceView* pViewNULL = NULL; - m_dx11Context->CSSetShaderResources( 0, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 1, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 2, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 3, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 4, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 5, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 6, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 7, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 8, 1, &pViewNULL ); - - ID3D11UnorderedAccessView* pUAViewNULL = NULL; - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &pUAViewNULL, NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &pUAViewNULL, NULL ); - - ID3D11Buffer *pBufferNull = NULL; - m_dx11Context->CSSetConstantBuffers( 0, 1, &pBufferNull ); - } - - -} // btDX11SIMDAwareSoftBodySolver::applyForces - -/** - * Integrate motion on the solver. - */ -void btDX11SIMDAwareSoftBodySolver::integrate( float solverdt ) -{ - // TEMPORARY COPIES + m_linkData.moveToAccelerator(); m_vertexData.moveToAccelerator(); - // No need to batch link solver, it is entirely parallel - // Copy kernel parameters to GPU - IntegrateCB constBuffer; + - constBuffer.numNodes = m_vertexData.getNumVertices(); - constBuffer.solverdt = solverdt; - - // Todo: factor this out. Number of nodes is static and sdt might be, too, we can update this just once on setup - D3D11_MAPPED_SUBRESOURCE MappedResource = {0}; - m_dx11Context->Map( integrateKernel.constBuffer, 0, D3D11_MAP_WRITE_DISCARD, 0, &MappedResource ); - memcpy( MappedResource.pData, &constBuffer, sizeof(IntegrateCB) ); - m_dx11Context->Unmap( integrateKernel.constBuffer, 0 ); - m_dx11Context->CSSetConstantBuffers( 0, 1, &integrateKernel.constBuffer ); + prepareCollisionConstraints(); - // Set resources and dispatch - m_dx11Context->CSSetShaderResources( 0, 1, &(m_vertexData.m_dx11VertexInverseMass.getSRV()) ); - - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &(m_vertexData.m_dx11VertexPosition.getUAV()), NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &(m_vertexData.m_dx11VertexVelocity.getUAV()), NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 2, 1, &(m_vertexData.m_dx11VertexPreviousPosition.getUAV()), NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 3, 1, &(m_vertexData.m_dx11VertexForceAccumulator.getUAV()), NULL ); - - // Execute the kernel - m_dx11Context->CSSetShader( integrateKernel.kernel, NULL, 0 ); - - int numBlocks = (constBuffer.numNodes + (128-1)) / 128; - m_dx11Context->Dispatch(numBlocks, 1, 1 ); + // Solve drift + for( int iteration = 0; iteration < m_numberOfPositionIterations ; ++iteration ) { - // Tidy up - ID3D11ShaderResourceView* pViewNULL = NULL; - m_dx11Context->CSSetShaderResources( 0, 1, &pViewNULL ); - ID3D11UnorderedAccessView* pUAViewNULL = NULL; - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &pUAViewNULL, NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &pUAViewNULL, NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 2, 1, &pUAViewNULL, NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 3, 1, &pUAViewNULL, NULL ); + for( int i = 0; i < m_linkData.m_wavefrontBatchStartLengths.size(); ++i ) + { + int startWave = m_linkData.m_wavefrontBatchStartLengths[i].start; + int numWaves = m_linkData.m_wavefrontBatchStartLengths[i].length; - ID3D11Buffer *pBufferNull = NULL; - m_dx11Context->CSSetConstantBuffers( 0, 1, &pBufferNull ); - } -} // btDX11SIMDAwareSoftBodySolver::integrate + solveLinksForPosition( startWave, numWaves, kst, ti ); + } + + } // for( int iteration = 0; iteration < m_numberOfPositionIterations ; ++iteration ) + + + + + // At this point assume that the force array is blank - we will overwrite it + solveCollisionsAndUpdateVelocities( 1.f/solverdt ); + +} // btDX11SIMDAwareSoftBodySolver::solveConstraints -float btDX11SIMDAwareSoftBodySolver::computeTriangleArea( - const Vectormath::Aos::Point3 &vertex0, - const Vectormath::Aos::Point3 &vertex1, - const Vectormath::Aos::Point3 &vertex2 ) -{ - Vectormath::Aos::Vector3 a = vertex1 - vertex0; - Vectormath::Aos::Vector3 b = vertex2 - vertex0; - Vectormath::Aos::Vector3 crossProduct = cross(a, b); - float area = length( crossProduct ); - return area; -} // btDX11SIMDAwareSoftBodySolver::computeTriangleArea -// Update constants here is a simple CPU version that is run on optimize void btDX11SIMDAwareSoftBodySolver::updateConstants( float timeStep ) { using namespace Vectormath::Aos; @@ -715,106 +408,10 @@ void btDX11SIMDAwareSoftBodySolver::updateConstants( float timeStep ) } } // btDX11SIMDAwareSoftBodySolver::updateConstants - - -void btDX11SIMDAwareSoftBodySolver::solveConstraints( float solverdt ) -{ - - //std::cerr << "'GPU' solve constraints\n"; - using Vectormath::Aos::Vector3; - using Vectormath::Aos::Point3; - using Vectormath::Aos::lengthSqr; - using Vectormath::Aos::dot; - - // Prepare links - int numLinks = m_linkData.getNumLinks(); - int numVertices = m_vertexData.getNumVertices(); - - float kst = 1.f; - float ti = 0.f; - - - m_dx11PerClothDampingFactor.moveToGPU(); - m_dx11PerClothVelocityCorrectionCoefficient.moveToGPU(); - - - - // Ensure data is on accelerator - m_linkData.moveToAccelerator(); - m_vertexData.moveToAccelerator(); - - // Solve drift - for( int iteration = 0; iteration < m_numberOfPositionIterations ; ++iteration ) - { - int it = iteration; - - for( int i = 0; i < m_linkData.m_wavefrontBatchStartLengths.size(); ++i ) - { - int startWave = m_linkData.m_wavefrontBatchStartLengths[i].start; - int numWaves = m_linkData.m_wavefrontBatchStartLengths[i].length; - - solveLinksForPosition( startWave, numWaves, kst, ti ); - } - - } // for( int iteration = 0; iteration < m_numberOfPositionIterations ; ++iteration ) - - - - - updateVelocitiesFromPositionsWithoutVelocities( 1.f/solverdt ); - -} // btDX11SIMDAwareSoftBodySolver::solveConstraints - - - - ////////////////////////////////////// // Kernel dispatches -void btDX11SIMDAwareSoftBodySolver::updatePositionsFromVelocities( float solverdt ) -{ - // No need to batch link solver, it is entirely parallel - // Copy kernel parameters to GPU - UpdatePositionsFromVelocitiesCB constBuffer; - - constBuffer.numNodes = m_vertexData.getNumVertices(); - constBuffer.solverSDT = solverdt; - - // Todo: factor this out. Number of nodes is static and sdt might be, too, we can update this just once on setup - D3D11_MAPPED_SUBRESOURCE MappedResource = {0}; - m_dx11Context->Map( updatePositionsFromVelocitiesKernel.constBuffer, 0, D3D11_MAP_WRITE_DISCARD, 0, &MappedResource ); - memcpy( MappedResource.pData, &constBuffer, sizeof(UpdatePositionsFromVelocitiesCB) ); - m_dx11Context->Unmap( updatePositionsFromVelocitiesKernel.constBuffer, 0 ); - m_dx11Context->CSSetConstantBuffers( 0, 1, &updatePositionsFromVelocitiesKernel.constBuffer ); - - // Set resources and dispatch - m_dx11Context->CSSetShaderResources( 0, 1, &(m_vertexData.m_dx11VertexVelocity.getSRV()) ); - - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &(m_vertexData.m_dx11VertexPreviousPosition.getUAV()), NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &(m_vertexData.m_dx11VertexPosition.getUAV()), NULL ); - - // Execute the kernel - m_dx11Context->CSSetShader( updatePositionsFromVelocitiesKernel.kernel, NULL, 0 ); - - int numBlocks = (constBuffer.numNodes + (128-1)) / 128; - m_dx11Context->Dispatch(numBlocks, 1, 1 ); - - { - // Tidy up - ID3D11ShaderResourceView* pViewNULL = NULL; - m_dx11Context->CSSetShaderResources( 0, 1, &pViewNULL ); - - ID3D11UnorderedAccessView* pUAViewNULL = NULL; - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &pUAViewNULL, NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &pUAViewNULL, NULL ); - - ID3D11Buffer *pBufferNull = NULL; - m_dx11Context->CSSetConstantBuffers( 0, 1, &pBufferNull ); - } -} // btDX11SIMDAwareSoftBodySolver::updatePositionsFromVelocities - - void btDX11SIMDAwareSoftBodySolver::solveLinksForPosition( int startWave, int numWaves, float kst, float ti ) { @@ -873,107 +470,6 @@ void btDX11SIMDAwareSoftBodySolver::solveLinksForPosition( int startWave, int nu } // btDX11SIMDAwareSoftBodySolver::solveLinksForPosition -void btDX11SIMDAwareSoftBodySolver::updateVelocitiesFromPositionsWithVelocities( float isolverdt ) -{ - // Copy kernel parameters to GPU - UpdateVelocitiesFromPositionsWithVelocitiesCB constBuffer; - - // Set the first link of the batch - // and the batch size - constBuffer.numNodes = m_vertexData.getNumVertices(); - constBuffer.isolverdt = isolverdt; - - D3D11_MAPPED_SUBRESOURCE MappedResource = {0}; - m_dx11Context->Map( updateVelocitiesFromPositionsWithVelocitiesKernel.constBuffer, 0, D3D11_MAP_WRITE_DISCARD, 0, &MappedResource ); - memcpy( MappedResource.pData, &constBuffer, sizeof(UpdateVelocitiesFromPositionsWithVelocitiesCB) ); - m_dx11Context->Unmap( updateVelocitiesFromPositionsWithVelocitiesKernel.constBuffer, 0 ); - m_dx11Context->CSSetConstantBuffers( 0, 1, &updateVelocitiesFromPositionsWithVelocitiesKernel.constBuffer ); - - // Set resources and dispatch - m_dx11Context->CSSetShaderResources( 0, 1, &(m_vertexData.m_dx11VertexPosition.getSRV()) ); - m_dx11Context->CSSetShaderResources( 1, 1, &(m_vertexData.m_dx11VertexPreviousPosition.getSRV()) ); - m_dx11Context->CSSetShaderResources( 2, 1, &(m_vertexData.m_dx11ClothIdentifier.getSRV()) ); - m_dx11Context->CSSetShaderResources( 3, 1, &(m_dx11PerClothVelocityCorrectionCoefficient.getSRV()) ); - m_dx11Context->CSSetShaderResources( 4, 1, &(m_dx11PerClothDampingFactor.getSRV()) ); - - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &(m_vertexData.m_dx11VertexVelocity.getUAV()), NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &(m_vertexData.m_dx11VertexForceAccumulator.getUAV()), NULL ); - - - // Execute the kernel - m_dx11Context->CSSetShader( updateVelocitiesFromPositionsWithVelocitiesKernel.kernel, NULL, 0 ); - - int numBlocks = (constBuffer.numNodes + (128-1)) / 128; - m_dx11Context->Dispatch(numBlocks , 1, 1 ); - - { - // Tidy up - ID3D11ShaderResourceView* pViewNULL = NULL; - m_dx11Context->CSSetShaderResources( 0, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 1, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 2, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 3, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 4, 1, &pViewNULL ); - - ID3D11UnorderedAccessView* pUAViewNULL = NULL; - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &pUAViewNULL, NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &pUAViewNULL, NULL ); - - ID3D11Buffer *pBufferNull = NULL; - m_dx11Context->CSSetConstantBuffers( 0, 1, &pBufferNull ); - } - -} // btDX11SIMDAwareSoftBodySolver::updateVelocitiesFromPositionsWithVelocities - -void btDX11SIMDAwareSoftBodySolver::updateVelocitiesFromPositionsWithoutVelocities( float isolverdt ) -{ - // Copy kernel parameters to GPU - UpdateVelocitiesFromPositionsWithoutVelocitiesCB constBuffer; - - // Set the first link of the batch - // and the batch size - constBuffer.numNodes = m_vertexData.getNumVertices(); - constBuffer.isolverdt = isolverdt; - - D3D11_MAPPED_SUBRESOURCE MappedResource = {0}; - m_dx11Context->Map( updateVelocitiesFromPositionsWithoutVelocitiesKernel.constBuffer, 0, D3D11_MAP_WRITE_DISCARD, 0, &MappedResource ); - memcpy( MappedResource.pData, &constBuffer, sizeof(UpdateVelocitiesFromPositionsWithoutVelocitiesCB) ); - m_dx11Context->Unmap( updateVelocitiesFromPositionsWithoutVelocitiesKernel.constBuffer, 0 ); - m_dx11Context->CSSetConstantBuffers( 0, 1, &updateVelocitiesFromPositionsWithoutVelocitiesKernel.constBuffer ); - - // Set resources and dispatch - m_dx11Context->CSSetShaderResources( 0, 1, &(m_vertexData.m_dx11VertexPosition.getSRV()) ); - m_dx11Context->CSSetShaderResources( 1, 1, &(m_vertexData.m_dx11VertexPreviousPosition.getSRV()) ); - m_dx11Context->CSSetShaderResources( 2, 1, &(m_vertexData.m_dx11ClothIdentifier.getSRV()) ); - m_dx11Context->CSSetShaderResources( 3, 1, &(m_dx11PerClothDampingFactor.getSRV()) ); - - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &(m_vertexData.m_dx11VertexVelocity.getUAV()), NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &(m_vertexData.m_dx11VertexForceAccumulator.getUAV()), NULL ); - - - // Execute the kernel - m_dx11Context->CSSetShader( updateVelocitiesFromPositionsWithoutVelocitiesKernel.kernel, NULL, 0 ); - - int numBlocks = (constBuffer.numNodes + (128-1)) / 128; - m_dx11Context->Dispatch(numBlocks , 1, 1 ); - - { - // Tidy up - ID3D11ShaderResourceView* pViewNULL = NULL; - m_dx11Context->CSSetShaderResources( 0, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 1, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 2, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 3, 1, &pViewNULL ); - - ID3D11UnorderedAccessView* pUAViewNULL = NULL; - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &pUAViewNULL, NULL ); - m_dx11Context->CSSetUnorderedAccessViews( 1, 1, &pUAViewNULL, NULL ); - - ID3D11Buffer *pBufferNull = NULL; - m_dx11Context->CSSetConstantBuffers( 0, 1, &pBufferNull ); - } - -} // btDX11SIMDAwareSoftBodySolver::updateVelocitiesFromPositionsWithoutVelocities // End kernel dispatches ///////////////////////////////////// @@ -986,213 +482,6 @@ void btDX11SIMDAwareSoftBodySolver::updateVelocitiesFromPositionsWithoutVelociti -btDX11SIMDAwareSoftBodySolver::btAcceleratedSoftBodyInterface *btDX11SIMDAwareSoftBodySolver::findSoftBodyInterface( const btSoftBody* const softBody ) -{ - for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) - { - btAcceleratedSoftBodyInterface *softBodyInterface = m_softBodySet[softBodyIndex]; - if( softBodyInterface->getSoftBody() == softBody ) - return softBodyInterface; - } - return 0; -} - -void btDX11SIMDAwareSoftBodySolver::copySoftBodyToVertexBuffer( const btSoftBody * const softBody, btVertexBufferDescriptor *vertexBuffer ) -{ - checkInitialized(); - - btAcceleratedSoftBodyInterface *currentCloth = findSoftBodyInterface( softBody ); - - const int firstVertex = currentCloth->getFirstVertex(); - const int lastVertex = firstVertex + currentCloth->getNumVertices(); - - if( vertexBuffer->getBufferType() == btVertexBufferDescriptor::CPU_BUFFER ) - { - // If we're doing a CPU-buffer copy must copy the data back to the host first - m_vertexData.m_dx11VertexPosition.copyFromGPU(); - m_vertexData.m_dx11VertexNormal.copyFromGPU(); - - const int firstVertex = currentCloth->getFirstVertex(); - const int lastVertex = firstVertex + currentCloth->getNumVertices(); - const btCPUVertexBufferDescriptor *cpuVertexBuffer = static_cast< btCPUVertexBufferDescriptor* >(vertexBuffer); - float *basePointer = cpuVertexBuffer->getBasePointer(); - - if( vertexBuffer->hasVertexPositions() ) - { - const int vertexOffset = cpuVertexBuffer->getVertexOffset(); - const int vertexStride = cpuVertexBuffer->getVertexStride(); - float *vertexPointer = basePointer + vertexOffset; - - for( int vertexIndex = firstVertex; vertexIndex < lastVertex; ++vertexIndex ) - { - Vectormath::Aos::Point3 position = m_vertexData.getPosition(vertexIndex); - *(vertexPointer + 0) = position.getX(); - *(vertexPointer + 1) = position.getY(); - *(vertexPointer + 2) = position.getZ(); - vertexPointer += vertexStride; - } - } - if( vertexBuffer->hasNormals() ) - { - const int normalOffset = cpuVertexBuffer->getNormalOffset(); - const int normalStride = cpuVertexBuffer->getNormalStride(); - float *normalPointer = basePointer + normalOffset; - - for( int vertexIndex = firstVertex; vertexIndex < lastVertex; ++vertexIndex ) - { - Vectormath::Aos::Vector3 normal = m_vertexData.getNormal(vertexIndex); - *(normalPointer + 0) = normal.getX(); - *(normalPointer + 1) = normal.getY(); - *(normalPointer + 2) = normal.getZ(); - normalPointer += normalStride; - } - } - } else if( vertexBuffer->getBufferType() == btVertexBufferDescriptor::DX11_BUFFER ) - { - // Do a DX11 copy shader DX to DX copy - - const btDX11VertexBufferDescriptor *dx11VertexBuffer = static_cast< btDX11VertexBufferDescriptor* >(vertexBuffer); - - // No need to batch link solver, it is entirely parallel - // Copy kernel parameters to GPU - OutputToVertexArrayCB constBuffer; - ID3D11ComputeShader* outputToVertexArrayShader = outputToVertexArrayWithoutNormalsKernel.kernel; - ID3D11Buffer* outputToVertexArrayConstBuffer = outputToVertexArrayWithoutNormalsKernel.constBuffer; - - constBuffer.startNode = firstVertex; - constBuffer.numNodes = currentCloth->getNumVertices(); - constBuffer.positionOffset = vertexBuffer->getVertexOffset(); - constBuffer.positionStride = vertexBuffer->getVertexStride(); - if( vertexBuffer->hasNormals() ) - { - constBuffer.normalOffset = vertexBuffer->getNormalOffset(); - constBuffer.normalStride = vertexBuffer->getNormalStride(); - outputToVertexArrayShader = outputToVertexArrayWithNormalsKernel.kernel; - outputToVertexArrayConstBuffer = outputToVertexArrayWithNormalsKernel.constBuffer; - } - - // TODO: factor this out. Number of nodes is static and sdt might be, too, we can update this just once on setup - D3D11_MAPPED_SUBRESOURCE MappedResource = {0}; - m_dx11Context->Map( outputToVertexArrayConstBuffer, 0, D3D11_MAP_WRITE_DISCARD, 0, &MappedResource ); - memcpy( MappedResource.pData, &constBuffer, sizeof(OutputToVertexArrayCB) ); - m_dx11Context->Unmap( outputToVertexArrayConstBuffer, 0 ); - m_dx11Context->CSSetConstantBuffers( 0, 1, &outputToVertexArrayConstBuffer ); - - // Set resources and dispatch - m_dx11Context->CSSetShaderResources( 0, 1, &(m_vertexData.m_dx11VertexPosition.getSRV()) ); - m_dx11Context->CSSetShaderResources( 1, 1, &(m_vertexData.m_dx11VertexNormal.getSRV()) ); - - ID3D11UnorderedAccessView* dx11UAV = dx11VertexBuffer->getDX11UAV(); - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &(dx11UAV), NULL ); - - // Execute the kernel - m_dx11Context->CSSetShader( outputToVertexArrayShader, NULL, 0 ); - - int numBlocks = (constBuffer.numNodes + (128-1)) / 128; - m_dx11Context->Dispatch(numBlocks, 1, 1 ); - - { - // Tidy up - ID3D11ShaderResourceView* pViewNULL = NULL; - m_dx11Context->CSSetShaderResources( 0, 1, &pViewNULL ); - m_dx11Context->CSSetShaderResources( 1, 1, &pViewNULL ); - - ID3D11UnorderedAccessView* pUAViewNULL = NULL; - m_dx11Context->CSSetUnorderedAccessViews( 0, 1, &pUAViewNULL, NULL ); - - ID3D11Buffer *pBufferNull = NULL; - m_dx11Context->CSSetConstantBuffers( 0, 1, &pBufferNull ); - } - } - -} // btDX11SoftBodySolver::outputToVertexBuffers - - - - - -btDX11SIMDAwareSoftBodySolver::KernelDesc btDX11SIMDAwareSoftBodySolver::compileComputeShaderFromString( const char* shaderString, const char* shaderName, int constBufferSize, D3D10_SHADER_MACRO *compileMacros ) -{ - const char *cs5String = "cs_5_0"; - - HRESULT hr = S_OK; - ID3DBlob* pErrorBlob = NULL; - ID3DBlob* pBlob = NULL; - ID3D11ComputeShader* kernelPointer = 0; - - hr = D3DX11CompileFromMemory( - shaderString, - strlen(shaderString), - shaderName, - compileMacros, - NULL, - shaderName, - cs5String, - D3D10_SHADER_ENABLE_STRICTNESS, - NULL, - NULL, - &pBlob, - &pErrorBlob, - NULL - ); - - if( FAILED(hr) ) - { - if( pErrorBlob ) { - btAssert( "Compilation of compute shader failed\n" ); - char *debugString = (char*)pErrorBlob->GetBufferPointer(); - OutputDebugStringA( debugString ); - } - - SAFE_RELEASE( pErrorBlob ); - SAFE_RELEASE( pBlob ); - - btDX11SIMDAwareSoftBodySolver::KernelDesc descriptor; - descriptor.kernel = 0; - descriptor.constBuffer = 0; - return descriptor; - } - - // Create the Compute Shader - hr = m_dx11Device->CreateComputeShader( pBlob->GetBufferPointer(), pBlob->GetBufferSize(), NULL, &kernelPointer ); - if( FAILED( hr ) ) - { - btDX11SIMDAwareSoftBodySolver::KernelDesc descriptor; - descriptor.kernel = 0; - descriptor.constBuffer = 0; - return descriptor; - } - - ID3D11Buffer* constBuffer = 0; - if( constBufferSize > 0 ) - { - // Create the constant buffer - D3D11_BUFFER_DESC constant_buffer_desc; - ZeroMemory(&constant_buffer_desc, sizeof(constant_buffer_desc)); - constant_buffer_desc.ByteWidth = constBufferSize; - constant_buffer_desc.Usage = D3D11_USAGE_DYNAMIC; - constant_buffer_desc.BindFlags = D3D11_BIND_CONSTANT_BUFFER; - constant_buffer_desc.CPUAccessFlags = D3D11_CPU_ACCESS_WRITE; - m_dx11Device->CreateBuffer(&constant_buffer_desc, NULL, &constBuffer); - if( FAILED( hr ) ) - { - KernelDesc descriptor; - descriptor.kernel = 0; - descriptor.constBuffer = 0; - return descriptor; - } - } - - SAFE_RELEASE( pErrorBlob ); - SAFE_RELEASE( pBlob ); - - btDX11SIMDAwareSoftBodySolver::KernelDesc descriptor; - descriptor.kernel = kernelPointer; - descriptor.constBuffer = constBuffer; - return descriptor; -} // compileComputeShader - - bool btDX11SIMDAwareSoftBodySolver::buildShaders() { // Ensure current kernels are released first @@ -1205,7 +494,7 @@ bool btDX11SIMDAwareSoftBodySolver::buildShaders() return true; - updatePositionsFromVelocitiesKernel = compileComputeShaderFromString( UpdatePositionsFromVelocitiesHLSLString, "UpdatePositionsFromVelocitiesKernel", sizeof(UpdatePositionsFromVelocitiesCB) ); + updatePositionsFromVelocitiesKernel = dxFunctions.compileComputeShaderFromString( UpdatePositionsFromVelocitiesHLSLString, "UpdatePositionsFromVelocitiesKernel", sizeof(UpdatePositionsFromVelocitiesCB) ); if( !updatePositionsFromVelocitiesKernel.constBuffer ) returnVal = false; @@ -1223,67 +512,53 @@ bool btDX11SIMDAwareSoftBodySolver::buildShaders() D3D10_SHADER_MACRO solvePositionsMacros[6] = { "MAX_NUM_VERTICES_PER_WAVE", maxVerticesPerWavefront, "MAX_BATCHES_PER_WAVE", maxBatchesPerWavefront, "WAVEFRONT_SIZE", waveFrontSize, "WAVEFRONT_BLOCK_MULTIPLIER", waveFrontBlockMultiplier, "BLOCK_SIZE", blockSize, 0, 0 }; - solvePositionsFromLinksKernel = compileComputeShaderFromString( SolvePositionsSIMDBatchedHLSLString, "SolvePositionsFromLinksKernel", sizeof(SolvePositionsFromLinksKernelCB), solvePositionsMacros ); + solvePositionsFromLinksKernel = dxFunctions.compileComputeShaderFromString( SolvePositionsSIMDBatchedHLSLString, "SolvePositionsFromLinksKernel", sizeof(SolvePositionsFromLinksKernelCB), solvePositionsMacros ); if( !solvePositionsFromLinksKernel.constBuffer ) returnVal = false; - updateVelocitiesFromPositionsWithVelocitiesKernel = compileComputeShaderFromString( UpdateNodesHLSLString, "updateVelocitiesFromPositionsWithVelocitiesKernel", sizeof(UpdateVelocitiesFromPositionsWithVelocitiesCB) ); + updateVelocitiesFromPositionsWithVelocitiesKernel = dxFunctions.compileComputeShaderFromString( UpdateNodesHLSLString, "updateVelocitiesFromPositionsWithVelocitiesKernel", sizeof(UpdateVelocitiesFromPositionsWithVelocitiesCB) ); if( !updateVelocitiesFromPositionsWithVelocitiesKernel.constBuffer ) returnVal = false; - updateVelocitiesFromPositionsWithoutVelocitiesKernel = compileComputeShaderFromString( UpdatePositionsHLSLString, "updateVelocitiesFromPositionsWithoutVelocitiesKernel", sizeof(UpdateVelocitiesFromPositionsWithoutVelocitiesCB)); + updateVelocitiesFromPositionsWithoutVelocitiesKernel = dxFunctions.compileComputeShaderFromString( UpdatePositionsHLSLString, "updateVelocitiesFromPositionsWithoutVelocitiesKernel", sizeof(UpdateVelocitiesFromPositionsWithoutVelocitiesCB)); if( !updateVelocitiesFromPositionsWithoutVelocitiesKernel.constBuffer ) returnVal = false; - integrateKernel = compileComputeShaderFromString( IntegrateHLSLString, "IntegrateKernel", sizeof(IntegrateCB) ); + integrateKernel = dxFunctions.compileComputeShaderFromString( IntegrateHLSLString, "IntegrateKernel", sizeof(IntegrateCB) ); if( !integrateKernel.constBuffer ) returnVal = false; - applyForcesKernel = compileComputeShaderFromString( ApplyForcesHLSLString, "ApplyForcesKernel", sizeof(ApplyForcesCB) ); + applyForcesKernel = dxFunctions.compileComputeShaderFromString( ApplyForcesHLSLString, "ApplyForcesKernel", sizeof(ApplyForcesCB) ); if( !applyForcesKernel.constBuffer ) returnVal = false; - - // TODO: Rename to UpdateSoftBodies - resetNormalsAndAreasKernel = compileComputeShaderFromString( UpdateNormalsHLSLString, "ResetNormalsAndAreasKernel", sizeof(UpdateSoftBodiesCB) ); + solveCollisionsAndUpdateVelocitiesKernel = dxFunctions.compileComputeShaderFromString( SolveCollisionsAndUpdateVelocitiesHLSLString, "SolveCollisionsAndUpdateVelocitiesKernel", sizeof(SolveCollisionsAndUpdateVelocitiesCB) ); + if( !solveCollisionsAndUpdateVelocitiesKernel.constBuffer ) + returnVal = false; + resetNormalsAndAreasKernel = dxFunctions.compileComputeShaderFromString( UpdateNormalsHLSLString, "ResetNormalsAndAreasKernel", sizeof(UpdateSoftBodiesCB) ); if( !resetNormalsAndAreasKernel.constBuffer ) returnVal = false; - normalizeNormalsAndAreasKernel = compileComputeShaderFromString( UpdateNormalsHLSLString, "NormalizeNormalsAndAreasKernel", sizeof(UpdateSoftBodiesCB) ); + normalizeNormalsAndAreasKernel = dxFunctions.compileComputeShaderFromString( UpdateNormalsHLSLString, "NormalizeNormalsAndAreasKernel", sizeof(UpdateSoftBodiesCB) ); if( !normalizeNormalsAndAreasKernel.constBuffer ) returnVal = false; - updateSoftBodiesKernel = compileComputeShaderFromString( UpdateNormalsHLSLString, "UpdateSoftBodiesKernel", sizeof(UpdateSoftBodiesCB) ); + updateSoftBodiesKernel = dxFunctions.compileComputeShaderFromString( UpdateNormalsHLSLString, "UpdateSoftBodiesKernel", sizeof(UpdateSoftBodiesCB) ); if( !updateSoftBodiesKernel.constBuffer ) returnVal = false; - outputToVertexArrayWithNormalsKernel = compileComputeShaderFromString( OutputToVertexArrayHLSLString, "OutputToVertexArrayWithNormalsKernel", sizeof(OutputToVertexArrayCB) ); - if( !outputToVertexArrayWithNormalsKernel.constBuffer ) + + computeBoundsKernel = dxFunctions.compileComputeShaderFromString( ComputeBoundsHLSLString, "ComputeBoundsKernel", sizeof(ComputeBoundsCB) ); + if( !computeBoundsKernel.constBuffer ) returnVal = false; - outputToVertexArrayWithoutNormalsKernel = compileComputeShaderFromString( OutputToVertexArrayHLSLString, "OutputToVertexArrayWithoutNormalsKernel", sizeof(OutputToVertexArrayCB) ); - if( !outputToVertexArrayWithoutNormalsKernel.constBuffer ) - returnVal = false; - if( returnVal ) m_shadersInitialized = true; return returnVal; -} +} // btDX11SIMDAwareSoftBodySolver::buildShaders - - -void btDX11SIMDAwareSoftBodySolver::predictMotion( float timeStep ) +static Vectormath::Aos::Transform3 toTransform3( const btTransform &transform ) { - // Fill the force arrays with current acceleration data etc - m_perClothWindVelocity.resize( m_softBodySet.size() ); - for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) - { - btSoftBody *softBody = m_softBodySet[softBodyIndex]->getSoftBody(); - - m_perClothWindVelocity[softBodyIndex] = toVector3(softBody->getWindVelocity()); - } - m_dx11PerClothWindVelocity.changedOnCPU(); - - // Apply forces that we know about to the cloths - applyForces( timeStep * getTimeScale() ); - - // Itegrate motion for all soft bodies dealt with by the solver - integrate( timeStep * getTimeScale() ); - // End prediction work for solvers + Vectormath::Aos::Transform3 outTransform; + outTransform.setCol(0, toVector3(transform.getBasis().getColumn(0))); + outTransform.setCol(1, toVector3(transform.getBasis().getColumn(1))); + outTransform.setCol(2, toVector3(transform.getBasis().getColumn(2))); + outTransform.setCol(3, toVector3(transform.getOrigin())); + return outTransform; } @@ -1297,13 +572,6 @@ void btDX11SIMDAwareSoftBodySolver::predictMotion( float timeStep ) - - - - - - - static void generateBatchesOfWavefronts( btAlignedObjectArray < btAlignedObjectArray > &linksForWavefronts, btSoftBodyLinkData &linkData, int numVertices, btAlignedObjectArray < btAlignedObjectArray > &wavefrontBatches ) { // A per-batch map of truth values stating whether a given vertex is in that batch @@ -1406,9 +674,7 @@ template< typename T > static void insertUniqueAndOrderedIntoVector( btAlignedOb insertAtIndex( vectorToUpdate, index, element ); } -// Experimental batch generation that we could use in the simulations -// Attempts to generate larger batches that work on a per-wavefront basis -void generateLinksPerVertex( int numVertices, btSoftBodyLinkData &linkData, btAlignedObjectArray< int > &listOfLinksPerVertex, btAlignedObjectArray &numLinksPerVertex, int &maxLinks ) +static void generateLinksPerVertex( int numVertices, btSoftBodyLinkData &linkData, btAlignedObjectArray< int > &listOfLinksPerVertex, btAlignedObjectArray &numLinksPerVertex, int &maxLinks ) { for( int linkIndex = 0; linkIndex < linkData.getNumLinks(); ++linkIndex ) { @@ -1489,14 +755,6 @@ static void computeBatchingIntoWavefronts( generateLinksPerVertex( numVertices, linkData, listOfLinksPerVertex, numLinksPerVertex, maxLinksPerVertex ); - for( int vertex = 0; vertex < 10; ++vertex ) - { - for( int link = 0; link < numLinksPerVertex[vertex]; ++link ) - { - int linkAddress = vertex * maxLinksPerVertex + link; - } - } - // At this point we know what links we have for each vertex so we can start batching diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11SIMDAware.h b/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11SIMDAware.h index ceac535e2..7b9cc2ce2 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11SIMDAware.h +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/DX11/btSoftBodySolver_DX11SIMDAware.h @@ -14,7 +14,7 @@ subject to the following restrictions: */ #include "vectormath/vmInclude.h" -#include "BulletSoftBody/btSoftBodySolvers.h" +#include "btSoftBodySolver_DX11.h" #include "btSoftBodySolverVertexBuffer_DX11.h" #include "btSoftBodySolverLinkData_DX11SIMDAware.h" #include "btSoftBodySolverVertexData_DX11.h" @@ -24,176 +24,9 @@ subject to the following restrictions: #ifndef BT_SOFT_BODY_DX11_SOLVER_SIMDAWARE_H #define BT_SOFT_BODY_DX11_SOLVER_SIMDAWARE_H -class btDX11SIMDAwareSoftBodySolver : public btSoftBodySolver +class btDX11SIMDAwareSoftBodySolver : public btDX11SoftBodySolver { -public: - - /** - * SoftBody class to maintain information about a soft body instance - * within a solver. - * This data addresses the main solver arrays. - */ - class btAcceleratedSoftBodyInterface - { - protected: - /** Current number of vertices that are part of this cloth */ - int m_numVertices; - /** Maximum number of vertices allocated to be part of this cloth */ - int m_maxVertices; - /** Current number of triangles that are part of this cloth */ - int m_numTriangles; - /** Maximum number of triangles allocated to be part of this cloth */ - int m_maxTriangles; - /** Index of first vertex in the world allocated to this cloth */ - int m_firstVertex; - /** Index of first triangle in the world allocated to this cloth */ - int m_firstTriangle; - /** Index of first link in the world allocated to this cloth */ - int m_firstLink; - /** Maximum number of links allocated to this cloth */ - int m_maxLinks; - /** Current number of links allocated to this cloth */ - int m_numLinks; - - /** The actual soft body this data represents */ - btSoftBody *m_softBody; - - - public: - btAcceleratedSoftBodyInterface( btSoftBody *softBody ) : - m_softBody( softBody ) - { - m_numVertices = 0; - m_maxVertices = 0; - m_numTriangles = 0; - m_maxTriangles = 0; - m_firstVertex = 0; - m_firstTriangle = 0; - m_firstLink = 0; - m_maxLinks = 0; - m_numLinks = 0; - } - int getNumVertices() - { - return m_numVertices; - } - - int getNumTriangles() - { - return m_numTriangles; - } - - int getMaxVertices() - { - return m_maxVertices; - } - - int getMaxTriangles() - { - return m_maxTriangles; - } - - int getFirstVertex() - { - return m_firstVertex; - } - - int getFirstTriangle() - { - return m_firstTriangle; - } - - - void setNumVertices( int numVertices ) - { - m_numVertices = numVertices; - } - - void setNumTriangles( int numTriangles ) - { - m_numTriangles = numTriangles; - } - - void setMaxVertices( int maxVertices ) - { - m_maxVertices = maxVertices; - } - - void setMaxTriangles( int maxTriangles ) - { - m_maxTriangles = maxTriangles; - } - - void setFirstVertex( int firstVertex ) - { - m_firstVertex = firstVertex; - } - - void setFirstTriangle( int firstTriangle ) - { - m_firstTriangle = firstTriangle; - } - - void setMaxLinks( int maxLinks ) - { - m_maxLinks = maxLinks; - } - - void setNumLinks( int numLinks ) - { - m_numLinks = numLinks; - } - - void setFirstLink( int firstLink ) - { - m_firstLink = firstLink; - } - - int getMaxLinks() - { - return m_maxLinks; - } - - int getNumLinks() - { - return m_numLinks; - } - - int getFirstLink() - { - return m_firstLink; - } - - btSoftBody* getSoftBody() - { - return m_softBody; - } - - }; - - - class KernelDesc - { - protected: - - - public: - ID3D11ComputeShader* kernel; - ID3D11Buffer* constBuffer; - - KernelDesc() - { - kernel = 0; - constBuffer = 0; - } - - virtual ~KernelDesc() - { - // TODO: this should probably destroy its kernel but we need to be careful - // in case KernelDescs are copied - } - }; - +protected: struct SolvePositionsFromLinksKernelCB { int startWave; @@ -202,230 +35,46 @@ public: float ti; }; - struct IntegrateCB - { - int numNodes; - float solverdt; - int padding1; - int padding2; - }; - - struct UpdatePositionsFromVelocitiesCB - { - int numNodes; - float solverSDT; - int padding1; - int padding2; - }; - - struct UpdateVelocitiesFromPositionsWithoutVelocitiesCB - { - int numNodes; - float isolverdt; - int padding1; - int padding2; - }; - - struct UpdateVelocitiesFromPositionsWithVelocitiesCB - { - int numNodes; - float isolverdt; - int padding1; - int padding2; - }; - - struct UpdateSoftBodiesCB - { - int numNodes; - int startFace; - int numFaces; - float epsilon; - }; - - - struct OutputToVertexArrayCB - { - int startNode; - int numNodes; - int positionOffset; - int positionStride; - - int normalOffset; - int normalStride; - int padding1; - int padding2; - }; - - - struct ApplyForcesCB - { - unsigned int numNodes; - float solverdt; - float epsilon; - int padding3; - }; - - struct AddVelocityCB - { - int startNode; - int lastNode; - float velocityX; - float velocityY; - float velocityZ; - int padding1; - int padding2; - int padding3; - }; - - -private: - ID3D11Device * m_dx11Device; - ID3D11DeviceContext* m_dx11Context; - /** Link data for all cloths. Note that this will be sorted batch-wise for efficient computation and m_linkAddresses will maintain the addressing. */ btSoftBodyLinkDataDX11SIMDAware m_linkData; - btSoftBodyVertexDataDX11 m_vertexData; - btSoftBodyTriangleDataDX11 m_triangleData; /** Variable to define whether we need to update solver constants on the next iteration */ bool m_updateSolverConstants; - bool m_shadersInitialized; - /** - * Cloths owned by this solver. - * Only our cloths are in this array. - */ - btAlignedObjectArray< btAcceleratedSoftBodyInterface * > m_softBodySet; + virtual bool buildShaders(); - /** Acceleration value to be applied to all non-static vertices in the solver. - * Index n is cloth n, array sized by number of cloths in the world not the solver. - */ - btAlignedObjectArray< Vectormath::Aos::Vector3 > m_perClothAcceleration; - btDX11Buffer m_dx11PerClothAcceleration; - - /** Wind velocity to be applied normal to all non-static vertices in the solver. - * Index n is cloth n, array sized by number of cloths in the world not the solver. - */ - btAlignedObjectArray< Vectormath::Aos::Vector3 > m_perClothWindVelocity; - btDX11Buffer m_dx11PerClothWindVelocity; - - /** Velocity damping factor */ - btAlignedObjectArray< float > m_perClothDampingFactor; - btDX11Buffer m_dx11PerClothDampingFactor; - - /** Velocity correction coefficient */ - btAlignedObjectArray< float > m_perClothVelocityCorrectionCoefficient; - btDX11Buffer m_dx11PerClothVelocityCorrectionCoefficient; - - /** Lift parameter for wind effect on cloth. */ - btAlignedObjectArray< float > m_perClothLiftFactor; - btDX11Buffer m_dx11PerClothLiftFactor; - - /** Drag parameter for wind effect on cloth. */ - btAlignedObjectArray< float > m_perClothDragFactor; - btDX11Buffer m_dx11PerClothDragFactor; - - /** Density of the medium in which each cloth sits */ - btAlignedObjectArray< float > m_perClothMediumDensity; - btDX11Buffer m_dx11PerClothMediumDensity; - - KernelDesc solvePositionsFromLinksKernel; - KernelDesc integrateKernel; - KernelDesc addVelocityKernel; - KernelDesc updatePositionsFromVelocitiesKernel; - KernelDesc updateVelocitiesFromPositionsWithoutVelocitiesKernel; - KernelDesc updateVelocitiesFromPositionsWithVelocitiesKernel; - KernelDesc resetNormalsAndAreasKernel; - KernelDesc normalizeNormalsAndAreasKernel; - KernelDesc updateSoftBodiesKernel; - KernelDesc outputToVertexArrayWithNormalsKernel; - KernelDesc outputToVertexArrayWithoutNormalsKernel; - - KernelDesc outputToVertexArrayKernel; - KernelDesc applyForcesKernel; - KernelDesc collideSphereKernel; - KernelDesc collideCylinderKernel; - - - - /** - * Integrate motion on the solver. - */ - virtual void integrate( float solverdt ); - float computeTriangleArea( - const Vectormath::Aos::Point3 &vertex0, - const Vectormath::Aos::Point3 &vertex1, - const Vectormath::Aos::Point3 &vertex2 ); - - - /** - * Compile a compute shader kernel from a string and return the appropriate KernelDesc object. - */ - KernelDesc compileComputeShaderFromString( const char* shaderString, const char* shaderName, int constBufferSize, D3D10_SHADER_MACRO *compileMacros = 0 ); - - bool buildShaders(); - - void resetNormalsAndAreas( int numVertices ); - - void normalizeNormalsAndAreas( int numVertices ); - - void executeUpdateSoftBodies( int firstTriangle, int numTriangles ); - - Vectormath::Aos::Vector3 ProjectOnAxis( const Vectormath::Aos::Vector3 &v, const Vectormath::Aos::Vector3 &a ); - - void ApplyClampedForce( float solverdt, const Vectormath::Aos::Vector3 &force, const Vectormath::Aos::Vector3 &vertexVelocity, float inverseMass, Vectormath::Aos::Vector3 &vertexForce ); - - virtual void applyForces( float solverdt ); - void updateConstants( float timeStep ); - btAcceleratedSoftBodyInterface *findSoftBodyInterface( const btSoftBody* const softBody ); ////////////////////////////////////// // Kernel dispatches - void prepareLinks(); - - void updatePositionsFromVelocities( float solverdt ); - void solveLinksForPosition( int startLink, int numLinks, float kst, float ti ); - void solveLinksForVelocity( int startLink, int numLinks, float kst ); - void updateVelocitiesFromPositionsWithVelocities( float isolverdt ); - void updateVelocitiesFromPositionsWithoutVelocities( float isolverdt ); + + void solveLinksForPosition( int startLink, int numLinks, float kst, float ti ); // End kernel dispatches ///////////////////////////////////// - void releaseKernels(); public: btDX11SIMDAwareSoftBodySolver(ID3D11Device * dx11Device, ID3D11DeviceContext* dx11Context); virtual ~btDX11SIMDAwareSoftBodySolver(); - - virtual btSoftBodyLinkData &getLinkData(); - virtual btSoftBodyVertexData &getVertexData(); - - virtual btSoftBodyTriangleData &getTriangleData(); - - - - virtual bool checkInitialized(); - - virtual void updateSoftBodies( ); - virtual void optimize( btAlignedObjectArray< btSoftBody * > &softBodies ); virtual void solveConstraints( float solverdt ); - - virtual void predictMotion( float solverdt ); - - virtual void copySoftBodyToVertexBuffer( const btSoftBody *const softBody, btVertexBufferDescriptor *vertexBuffer ); + + virtual SolverTypes getSolverType() const + { + return DX_SIMD_SOLVER; + } + }; #endif // #ifndef BT_SOFT_BODY_DX11_SOLVER_SIMDAWARE_H diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/AMD/CMakeLists.txt b/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/AMD/CMakeLists.txt index 10cd4dfe4..38e5a6620 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/AMD/CMakeLists.txt +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/AMD/CMakeLists.txt @@ -11,15 +11,21 @@ ADD_DEFINITIONS(-DCL_PLATFORM_AMD) SET(BulletSoftBodyOpenCLSolvers_SRCS ../btSoftBodySolver_OpenCL.cpp + ../btSoftBodySolver_OpenCLSIMDAware.cpp + ../btSoftBodySolverOutputCLtoGL.cpp ) SET(BulletSoftBodyOpenCLSolvers_HDRS ../btSoftBodySolver_OpenCL.h + ../btSoftBodySolver_OpenCLSIMDAware.h ../../CPU/btSoftBodySolverData.h ../btSoftBodySolverVertexData_OpenCL.h ../btSoftBodySolverTriangleData_OpenCL.h ../btSoftBodySolverLinkData_OpenCL.h + ../btSoftBodySolverLinkData_OpenCLSIMDAware.h ../btSoftBodySolverBuffer_OpenCL.h + ../btSoftBodySolverVertexBuffer_OpenGL.h + ../btSoftBodySolverOutputCLtoGL.h ) # OpenCL and HLSL Shaders. diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/NVidia/CMakeLists.txt b/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/NVidia/CMakeLists.txt index ce38adfc7..f102d3693 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/NVidia/CMakeLists.txt +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/NVidia/CMakeLists.txt @@ -5,8 +5,11 @@ INCLUDE_DIRECTORIES( ) + SET(BulletSoftBodyOpenCLSolvers_SRCS ../btSoftBodySolver_OpenCL.cpp + ../btSoftBodySolver_OpenCLSIMDAware.cpp + ../btSoftBodySolverOutputCLtoGL.cpp ) SET(BulletSoftBodyOpenCLSolvers_HDRS @@ -15,7 +18,10 @@ SET(BulletSoftBodyOpenCLSolvers_HDRS ../btSoftBodySolverVertexData_OpenCL.h ../btSoftBodySolverTriangleData_OpenCL.h ../btSoftBodySolverLinkData_OpenCL.h + ../btSoftBodySolverLinkData_OpenCLSIMDAware.h ../btSoftBodySolverBuffer_OpenCL.h + ../btSoftBodySolverVertexBuffer_OpenGL.h + ../btSoftBodySolverOutputCLtoGL.h ) # OpenCL and HLSL Shaders. diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolverBuffer_OpenCL.h b/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolverBuffer_OpenCL.h index 8fa58cd16..61d474a7d 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolverBuffer_OpenCL.h +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolverBuffer_OpenCL.h @@ -65,6 +65,9 @@ public: cl_mem_flags flags= m_readOnlyOnGPU ? CL_MEM_READ_ONLY : CL_MEM_READ_WRITE; size_t size = m_CPUBuffer->size() * sizeof(ElementType); + // At a minimum the buffer must exist + if( size == 0 ) + size = sizeof(ElementType); m_buffer = clCreateBuffer(m_clContext, flags, size, 0, &err); if( err != CL_SUCCESS ) { @@ -81,6 +84,7 @@ public: btOpenCLBuffer( cl_command_queue commandQue,cl_context ctx, btAlignedObjectArray< ElementType >* CPUBuffer, bool readOnly) :m_cqCommandQue(commandQue), m_clContext(ctx), + m_buffer(0), m_CPUBuffer(CPUBuffer), m_gpuSize(0), m_onGPU(false), @@ -91,6 +95,7 @@ public: ~btOpenCLBuffer() { + clReleaseMemObject(m_buffer); } @@ -105,6 +110,16 @@ public: m_onGPU = false; } + if( !m_allocated && m_CPUBuffer->size() == 0 ) + { + // If it isn't on the GPU and yet there is no data on the CPU side this may cause a problem with some kernels. + // We should create *something* on the device side + if (!createBuffer()) { + return false; + } + m_allocated = true; + } + if( !m_onGPU && m_CPUBuffer->size() > 0 ) { if (!m_allocated || (m_CPUBuffer->size() != m_gpuSize)) { diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCL.cpp b/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCL.cpp index 949cb8c24..ad2edfbda 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCL.cpp +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCL.cpp @@ -20,10 +20,26 @@ subject to the following restrictions: #include "btSoftBodySolver_OpenCL.h" #include "BulletSoftBody/btSoftBodySolverVertexBuffer.h" #include "BulletSoftBody/btSoftBody.h" +#include "BulletCollision/CollisionShapes/btCapsuleShape.h" +#include "LinearMath/btQuickprof.h" + +#ifdef USE_MINICL + #include "MiniCL/cl.h" +#else //USE_MINICL + #ifdef __APPLE__ + #include + #else + #include + #endif //__APPLE__ +#endif//USE_MINICL #define BT_DEFAULT_WORKGROUPSIZE 128 + +#define RELEASE_CL_KERNEL(kernelName) {if( kernelName ){ clReleaseKernel( kernelName ); kernelName = 0; }} + + //CL_VERSION_1_1 seems broken on NVidia SDK so just disable it #if (0)//CL_VERSION_1_1 == 1) @@ -49,6 +65,10 @@ static char* UpdateNormalsCLString = #include "OpenCLC/UpdateNormals.cl" static char* VSolveLinksCLString = #include "OpenCLC/VSolveLinks.cl" +static char* ComputeBoundsCLString = +#include "OpenCLC/ComputeBounds.cl" +static char* SolveCollisionsAndUpdateVelocitiesCLString = +#include "OpenCLC/SolveCollisionsAndUpdateVelocities.cl" #else ////OpenCL 1.0 kernels don't use float3 #define MSTRINGIFY(A) #A @@ -72,6 +92,10 @@ static char* UpdateNormalsCLString = #include "OpenCLC10/UpdateNormals.cl" static char* VSolveLinksCLString = #include "OpenCLC10/VSolveLinks.cl" +static char* ComputeBoundsCLString = +#include "OpenCLC10/ComputeBounds.cl" +static char* SolveCollisionsAndUpdateVelocitiesCLString = +#include "OpenCLC10/SolveCollisionsAndUpdateVelocities.cl" #endif //CL_VERSION_1_1 @@ -583,6 +607,7 @@ btOpenCLSoftBodySolver::btOpenCLSoftBodySolver(cl_command_queue queue, cl_contex m_linkData(queue, ctx), m_vertexData(queue, ctx), m_triangleData(queue, ctx), + clFunctions(queue, ctx), m_clPerClothAcceleration(queue, ctx, &m_perClothAcceleration, true ), m_clPerClothWindVelocity(queue, ctx, &m_perClothWindVelocity, true ), m_clPerClothDampingFactor(queue,ctx, &m_perClothDampingFactor, true ), @@ -590,6 +615,11 @@ btOpenCLSoftBodySolver::btOpenCLSoftBodySolver(cl_command_queue queue, cl_contex m_clPerClothLiftFactor(queue, ctx,&m_perClothLiftFactor, true ), m_clPerClothDragFactor(queue, ctx,&m_perClothDragFactor, true ), m_clPerClothMediumDensity(queue, ctx,&m_perClothMediumDensity, true ), + m_clPerClothCollisionObjects( queue, ctx, &m_perClothCollisionObjects, true ), + m_clCollisionObjectDetails( queue, ctx, &m_collisionObjectDetails, true ), + m_clPerClothMinBounds( queue, ctx, &m_perClothMinBounds, false ), + m_clPerClothMaxBounds( queue, ctx, &m_perClothMaxBounds, false ), + m_clPerClothFriction( queue, ctx, &m_perClothFriction, false ), m_cqCommandQue( queue ), m_cxMainContext(ctx), m_defaultWorkGroupSize(BT_DEFAULT_WORKGROUPSIZE) @@ -600,15 +630,85 @@ btOpenCLSoftBodySolver::btOpenCLSoftBodySolver(cl_command_queue queue, cl_contex m_updateSolverConstants = true; m_shadersInitialized = false; + + prepareLinksKernel = 0; + solvePositionsFromLinksKernel = 0; + updateConstantsKernel = 0; + integrateKernel = 0; + addVelocityKernel = 0; + updatePositionsFromVelocitiesKernel = 0; + updateVelocitiesFromPositionsWithoutVelocitiesKernel = 0; + updateVelocitiesFromPositionsWithVelocitiesKernel = 0; + vSolveLinksKernel = 0; + solveCollisionsAndUpdateVelocitiesKernel = 0; + resetNormalsAndAreasKernel = 0; + resetNormalsAndAreasKernel = 0; + normalizeNormalsAndAreasKernel = 0; + computeBoundsKernel = 0; + outputToVertexArrayKernel = 0; + applyForcesKernel = 0; } btOpenCLSoftBodySolver::~btOpenCLSoftBodySolver() { + releaseKernels(); } -void btOpenCLSoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * > &softBodies ) +void btOpenCLSoftBodySolver::releaseKernels() { - if( m_softBodySet.size() != softBodies.size() ) + RELEASE_CL_KERNEL( prepareLinksKernel ); + RELEASE_CL_KERNEL( solvePositionsFromLinksKernel ); + RELEASE_CL_KERNEL( updateConstantsKernel ); + RELEASE_CL_KERNEL( integrateKernel ); + RELEASE_CL_KERNEL( addVelocityKernel ); + RELEASE_CL_KERNEL( updatePositionsFromVelocitiesKernel ); + RELEASE_CL_KERNEL( updateVelocitiesFromPositionsWithoutVelocitiesKernel ); + RELEASE_CL_KERNEL( updateVelocitiesFromPositionsWithVelocitiesKernel ); + RELEASE_CL_KERNEL( vSolveLinksKernel ); + RELEASE_CL_KERNEL( solveCollisionsAndUpdateVelocitiesKernel ); + RELEASE_CL_KERNEL( resetNormalsAndAreasKernel ); + RELEASE_CL_KERNEL( normalizeNormalsAndAreasKernel ); + RELEASE_CL_KERNEL( computeBoundsKernel ); + RELEASE_CL_KERNEL( outputToVertexArrayKernel ); + RELEASE_CL_KERNEL( applyForcesKernel ); + + m_shadersInitialized = false; +} + +void btOpenCLSoftBodySolver::copyBackToSoftBodies() +{ + // Move the vertex data back to the host first + m_vertexData.moveFromAccelerator(); + + // Loop over soft bodies, copying all the vertex positions back for each body in turn + for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) + { + btOpenCLAcceleratedSoftBodyInterface *softBodyInterface = m_softBodySet[ softBodyIndex ]; + btSoftBody *softBody = softBodyInterface->getSoftBody(); + + int firstVertex = softBodyInterface->getFirstVertex(); + int numVertices = softBodyInterface->getNumVertices(); + + // Copy vertices from solver back into the softbody + for( int vertex = 0; vertex < numVertices; ++vertex ) + { + using Vectormath::Aos::Point3; + Point3 vertexPosition( getVertexData().getVertexPositions()[firstVertex + vertex] ); + + softBody->m_nodes[vertex].m_x.setX( vertexPosition.getX() ); + softBody->m_nodes[vertex].m_x.setY( vertexPosition.getY() ); + softBody->m_nodes[vertex].m_x.setZ( vertexPosition.getZ() ); + + softBody->m_nodes[vertex].m_n.setX( vertexPosition.getX() ); + softBody->m_nodes[vertex].m_n.setY( vertexPosition.getY() ); + softBody->m_nodes[vertex].m_n.setZ( vertexPosition.getZ() ); + } + } +} // btOpenCLSoftBodySolver::copyBackToSoftBodies + +void btOpenCLSoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * > &softBodies, bool forceUpdate ) +{ + if( forceUpdate || m_softBodySet.size() != softBodies.size() ) { // Have a change in the soft body set so update, reloading all the data getVertexData().clear(); @@ -633,6 +733,11 @@ void btOpenCLSoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * > &sof m_perClothLiftFactor.push_back( softBody->m_cfg.kLF ); m_perClothDragFactor.push_back( softBody->m_cfg.kDG ); m_perClothMediumDensity.push_back(softBody->getWorldInfo()->air_density); + // Simple init values. Actually we'll put 0 and -1 into them at the appropriate time + m_perClothMinBounds.push_back( UIntVector3(UINT_MAX, UINT_MAX, UINT_MAX) ); + m_perClothMaxBounds.push_back( UIntVector3(0, 0, 0) ); + m_perClothFriction.push_back( softBody->getFriction() ); + m_perClothCollisionObjects.push_back( CollisionObjectIndices(-1, -1) ); // Add space for new vertices and triangles in the default solver for now // TODO: Include space here for tearing too later @@ -738,12 +843,6 @@ btSoftBodyTriangleData &btOpenCLSoftBodySolver::getTriangleData() return m_triangleData; } - -bool btOpenCLSoftBodySolver::checkInitialized() -{ - return buildShaders(); -} - void btOpenCLSoftBodySolver::resetNormalsAndAreas( int numVertices ) { cl_int ciErrNum; @@ -751,11 +850,15 @@ void btOpenCLSoftBodySolver::resetNormalsAndAreas( int numVertices ) ciErrNum = clSetKernelArg(resetNormalsAndAreasKernel, 1, sizeof(cl_mem), (void*)&m_vertexData.m_clVertexNormal.m_buffer);//oclCHECKERROR(ciErrNum, CL_SUCCESS); ciErrNum = clSetKernelArg(resetNormalsAndAreasKernel, 2, sizeof(cl_mem), (void*)&m_vertexData.m_clVertexArea.m_buffer); //oclCHECKERROR(ciErrNum, CL_SUCCESS); size_t numWorkItems = m_defaultWorkGroupSize*((numVertices + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize); - ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue, resetNormalsAndAreasKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize, 0,0,0 ); - if( ciErrNum != CL_SUCCESS ) + if (numWorkItems) { - btAssert( 0 && "enqueueNDRangeKernel(resetNormalsAndAreasKernel)" ); + ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue, resetNormalsAndAreasKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize, 0,0,0 ); + + if( ciErrNum != CL_SUCCESS ) + { + btAssert( 0 && "enqueueNDRangeKernel(resetNormalsAndAreasKernel)" ); + } } } @@ -770,10 +873,13 @@ void btOpenCLSoftBodySolver::normalizeNormalsAndAreas( int numVertices ) ciErrNum = clSetKernelArg(normalizeNormalsAndAreasKernel, 2, sizeof(cl_mem), &m_vertexData.m_clVertexNormal.m_buffer); ciErrNum = clSetKernelArg(normalizeNormalsAndAreasKernel, 3, sizeof(cl_mem), &m_vertexData.m_clVertexArea.m_buffer); size_t numWorkItems = m_defaultWorkGroupSize*((numVertices + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize); - ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue, normalizeNormalsAndAreasKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize, 0,0,0); - if( ciErrNum != CL_SUCCESS ) + if (numWorkItems) { - btAssert( 0 && "enqueueNDRangeKernel(normalizeNormalsAndAreasKernel)"); + ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue, normalizeNormalsAndAreasKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize, 0,0,0); + if( ciErrNum != CL_SUCCESS ) + { + btAssert( 0 && "enqueueNDRangeKernel(normalizeNormalsAndAreasKernel)"); + } } } @@ -875,10 +981,13 @@ void btOpenCLSoftBodySolver::applyForces( float solverdt ) ciErrNum = clSetKernelArg(applyForcesKernel,12, sizeof(cl_mem), &m_vertexData.m_clVertexForceAccumulator.m_buffer); ciErrNum = clSetKernelArg(applyForcesKernel,13, sizeof(cl_mem), &m_vertexData.m_clVertexVelocity.m_buffer); size_t numWorkItems = m_defaultWorkGroupSize*((m_vertexData.getNumVertices() + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize); - ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,applyForcesKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize, 0,0,0); - if( ciErrNum != CL_SUCCESS ) + if (numWorkItems) { - btAssert( 0 && "enqueueNDRangeKernel(applyForcesKernel)"); + ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,applyForcesKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize, 0,0,0); + if( ciErrNum != CL_SUCCESS ) + { + btAssert( 0 && "enqueueNDRangeKernel(applyForcesKernel)"); + } } } @@ -904,10 +1013,13 @@ void btOpenCLSoftBodySolver::integrate( float solverdt ) ciErrNum = clSetKernelArg(integrateKernel, 6, sizeof(cl_mem), &m_vertexData.m_clVertexForceAccumulator.m_buffer); size_t numWorkItems = m_defaultWorkGroupSize*((m_vertexData.getNumVertices() + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize); - ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,integrateKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize,0,0,0); - if( ciErrNum != CL_SUCCESS ) + if (numWorkItems) { - btAssert( 0 && "enqueueNDRangeKernel(integrateKernel)"); + ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,integrateKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize,0,0,0); + if( ciErrNum != CL_SUCCESS ) + { + btAssert( 0 && "enqueueNDRangeKernel(integrateKernel)"); + } } } @@ -924,6 +1036,102 @@ float btOpenCLSoftBodySolver::computeTriangleArea( return area; } + +void btOpenCLSoftBodySolver::updateBounds() +{ + +//#define USE_GPU_BOUNDS_COMPUTATION +#ifdef USE_GPU_BOUNDS_COMPUTATION + using Vectormath::Aos::Point3; + // Interpretation structure for float and int + + struct FPRep { + unsigned int mantissa : 23; + unsigned int exponent : 8; + unsigned int sign : 1; + }; + union FloatAsInt + { + float floatValue; + int intValue; + unsigned int uintValue; + FPRep fpRep; + }; + + + // Update bounds array to min and max int values to allow easy atomics + for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) + { + m_perClothMinBounds[softBodyIndex] = UIntVector3( UINT_MAX, UINT_MAX, UINT_MAX ); + m_perClothMaxBounds[softBodyIndex] = UIntVector3( 0, 0, 0 ); + } + + m_vertexData.moveToAccelerator(); + m_clPerClothMinBounds.moveToGPU(); + m_clPerClothMaxBounds.moveToGPU(); + + + computeBounds( ); + + + m_clPerClothMinBounds.moveFromGPU(); + m_clPerClothMaxBounds.moveFromGPU(); + + + + for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) + { + UIntVector3 minBoundUInt = m_perClothMinBounds[softBodyIndex]; + UIntVector3 maxBoundUInt = m_perClothMaxBounds[softBodyIndex]; + /*UIntVector3 minBoundUInt; + minBoundUInt.x = m_perClothMinBounds[softBodyIndex*4]; + minBoundUInt.y = m_perClothMinBounds[softBodyIndex*4+1]; + minBoundUInt.z = m_perClothMinBounds[softBodyIndex*4+2]; + UIntVector3 maxBoundUInt; + maxBoundUInt.x = m_perClothMaxBounds[softBodyIndex*4]; + maxBoundUInt.y = m_perClothMaxBounds[softBodyIndex*4+1]; + maxBoundUInt.z = m_perClothMaxBounds[softBodyIndex*4+2];*/ + + // Convert back to float + FloatAsInt fai; + + btVector3 minBound; + fai.uintValue = minBoundUInt.x; + fai.uintValue ^= (((fai.uintValue >> 31) - 1) | 0x80000000); + minBound.setX( fai.floatValue ); + fai.uintValue = minBoundUInt.y; + fai.uintValue ^= (((fai.uintValue >> 31) - 1) | 0x80000000); + minBound.setY( fai.floatValue ); + fai.uintValue = minBoundUInt.z; + fai.uintValue ^= (((fai.uintValue >> 31) - 1) | 0x80000000); + minBound.setZ( fai.floatValue ); + + btVector3 maxBound; + fai.uintValue = maxBoundUInt.x; + fai.uintValue ^= (((fai.uintValue >> 31) - 1) | 0x80000000); + maxBound.setX( fai.floatValue ); + fai.uintValue = maxBoundUInt.y; + fai.uintValue ^= (((fai.uintValue >> 31) - 1) | 0x80000000); + maxBound.setY( fai.floatValue ); + fai.uintValue = maxBoundUInt.z; + fai.uintValue ^= (((fai.uintValue >> 31) - 1) | 0x80000000); + maxBound.setZ( fai.floatValue ); + + + // And finally assign to the soft body + m_softBodySet[softBodyIndex]->updateBounds( minBound, maxBound ); + } +#else + for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) + { + btVector3 minBound(-1e30,-1e30,-1e30), maxBound(1e30,1e30,1e30); + m_softBodySet[softBodyIndex]->updateBounds( minBound, maxBound ); + } +#endif//USE_GPU_BOUNDS_COMPUTATION + +} // btOpenCLSoftBodySolver::updateBounds + + void btOpenCLSoftBodySolver::updateConstants( float timeStep ) { @@ -954,6 +1162,66 @@ void btOpenCLSoftBodySolver::updateConstants( float timeStep ) } +class QuickSortCompare +{ + public: + + bool operator() ( const CollisionShapeDescription& a, const CollisionShapeDescription& b ) + { + return ( a.softBodyIdentifier < b.softBodyIdentifier ); + } +}; + + +/** + * Sort the collision object details array and generate indexing into it for the per-cloth collision object array. + */ +void btOpenCLSoftBodySolver::prepareCollisionConstraints() +{ + // First do a simple sort on the collision objects + btAlignedObjectArray numObjectsPerClothPrefixSum; + btAlignedObjectArray numObjectsPerCloth; + numObjectsPerCloth.resize( m_softBodySet.size(), 0 ); + numObjectsPerClothPrefixSum.resize( m_softBodySet.size(), 0 ); + + + + m_collisionObjectDetails.quickSort( QuickSortCompare() ); + + if (!m_perClothCollisionObjects.size()) + return; + + // Generating indexing for perClothCollisionObjects + // First clear the previous values with the "no collision object for cloth" constant + for( int clothIndex = 0; clothIndex < m_perClothCollisionObjects.size(); ++clothIndex ) + { + m_perClothCollisionObjects[clothIndex].firstObject = -1; + m_perClothCollisionObjects[clothIndex].endObject = -1; + } + int currentCloth = 0; + int startIndex = 0; + for( int collisionObject = 0; collisionObject < m_collisionObjectDetails.size(); ++collisionObject ) + { + int nextCloth = m_collisionObjectDetails[collisionObject].softBodyIdentifier; + if( nextCloth != currentCloth ) + { + // Changed cloth in the array + // Set the end index and the range is what we need for currentCloth + m_perClothCollisionObjects[currentCloth].firstObject = startIndex; + m_perClothCollisionObjects[currentCloth].endObject = collisionObject; + currentCloth = nextCloth; + startIndex = collisionObject; + } + } + + // And update last cloth + m_perClothCollisionObjects[currentCloth].firstObject = startIndex; + m_perClothCollisionObjects[currentCloth].endObject = m_collisionObjectDetails.size(); + +} // btOpenCLSoftBodySolver::prepareCollisionConstraints + + + void btOpenCLSoftBodySolver::solveConstraints( float solverdt ) { @@ -993,6 +1261,9 @@ void btOpenCLSoftBodySolver::solveConstraints( float solverdt ) } } + + prepareCollisionConstraints(); + // Compute new positions from velocity // Also update the previous position so that our position computation is now based on the new position from the velocity solution // rather than based directly on the original positions @@ -1016,8 +1287,9 @@ void btOpenCLSoftBodySolver::solveConstraints( float solverdt ) } // for( int iteration = 0; iteration < m_numberOfPositionIterations ; ++iteration ) - - updateVelocitiesFromPositionsWithoutVelocities( 1.f/solverdt ); + + // At this point assume that the force array is blank - we will overwrite it + solveCollisionsAndUpdateVelocities( 1.f/solverdt ); } @@ -1158,19 +1430,88 @@ void btOpenCLSoftBodySolver::updateVelocitiesFromPositionsWithoutVelocities( flo } // updateVelocitiesFromPositionsWithoutVelocities + +void btOpenCLSoftBodySolver::computeBounds( ) +{ + m_vertexData.moveToAccelerator(); + + cl_int ciErrNum; + int numVerts = m_vertexData.getNumVertices(); + int numSoftBodies = m_softBodySet.size(); + ciErrNum = clSetKernelArg(computeBoundsKernel, 0, sizeof(int), &numVerts); + ciErrNum = clSetKernelArg(computeBoundsKernel, 1, sizeof(int), &numSoftBodies); + ciErrNum = clSetKernelArg(computeBoundsKernel, 2, sizeof(cl_mem),&m_vertexData.m_clClothIdentifier.m_buffer); + ciErrNum = clSetKernelArg(computeBoundsKernel, 3, sizeof(cl_mem),&m_vertexData.m_clVertexPosition.m_buffer); + ciErrNum = clSetKernelArg(computeBoundsKernel, 4, sizeof(cl_mem),&m_clPerClothMinBounds.m_buffer); + ciErrNum = clSetKernelArg(computeBoundsKernel, 5, sizeof(cl_mem),&m_clPerClothMaxBounds.m_buffer); + ciErrNum = clSetKernelArg(computeBoundsKernel, 6, sizeof(cl_uint4)*256,0); + ciErrNum = clSetKernelArg(computeBoundsKernel, 7, sizeof(cl_uint4)*256,0); + + size_t numWorkItems = m_defaultWorkGroupSize*((m_vertexData.getNumVertices() + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize); + if (numWorkItems) + { + ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,computeBoundsKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize,0,0,0); + if( ciErrNum != CL_SUCCESS ) + { + btAssert( 0 && "enqueueNDRangeKernel(computeBoundsKernel)"); + } + } + clFinish(m_cqCommandQue); +} // btOpenCLSoftBodySolver::computeBounds + +void btOpenCLSoftBodySolver::solveCollisionsAndUpdateVelocities( float isolverdt ) +{ + + // Copy kernel parameters to GPU + m_vertexData.moveToAccelerator(); + m_clPerClothFriction.moveToGPU(); + m_clPerClothDampingFactor.moveToGPU(); + m_clPerClothCollisionObjects.moveToGPU(); + m_clCollisionObjectDetails.moveToGPU(); + + + cl_int ciErrNum; + int numVerts = m_vertexData.getNumVertices(); + ciErrNum = clSetKernelArg(solveCollisionsAndUpdateVelocitiesKernel, 0, sizeof(int), &numVerts); + ciErrNum = clSetKernelArg(solveCollisionsAndUpdateVelocitiesKernel, 1, sizeof(int), &isolverdt); + ciErrNum = clSetKernelArg(solveCollisionsAndUpdateVelocitiesKernel, 2, sizeof(cl_mem),&m_vertexData.m_clClothIdentifier.m_buffer); + ciErrNum = clSetKernelArg(solveCollisionsAndUpdateVelocitiesKernel, 3, sizeof(cl_mem),&m_vertexData.m_clVertexPreviousPosition.m_buffer); + ciErrNum = clSetKernelArg(solveCollisionsAndUpdateVelocitiesKernel, 4, sizeof(cl_mem),&m_clPerClothFriction.m_buffer); + ciErrNum = clSetKernelArg(solveCollisionsAndUpdateVelocitiesKernel, 5, sizeof(cl_mem),&m_clPerClothDampingFactor.m_buffer); + ciErrNum = clSetKernelArg(solveCollisionsAndUpdateVelocitiesKernel, 6, sizeof(cl_mem),&m_clPerClothCollisionObjects.m_buffer); + ciErrNum = clSetKernelArg(solveCollisionsAndUpdateVelocitiesKernel, 7, sizeof(cl_mem),&m_clCollisionObjectDetails.m_buffer); + ciErrNum = clSetKernelArg(solveCollisionsAndUpdateVelocitiesKernel, 8, sizeof(cl_mem),&m_vertexData.m_clVertexForceAccumulator.m_buffer); + ciErrNum = clSetKernelArg(solveCollisionsAndUpdateVelocitiesKernel, 9, sizeof(cl_mem),&m_vertexData.m_clVertexVelocity.m_buffer); + ciErrNum = clSetKernelArg(solveCollisionsAndUpdateVelocitiesKernel, 10, sizeof(cl_mem),&m_vertexData.m_clVertexPosition.m_buffer); + + size_t numWorkItems = m_defaultWorkGroupSize*((m_vertexData.getNumVertices() + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize); + if (numWorkItems) + { + ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,solveCollisionsAndUpdateVelocitiesKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize,0,0,0); + if( ciErrNum != CL_SUCCESS ) + { + btAssert( 0 && "enqueueNDRangeKernel(updateVelocitiesFromPositionsWithoutVelocitiesKernel)"); + } + } + +} // btOpenCLSoftBodySolver::updateVelocitiesFromPositionsWithoutVelocities + + + // End kernel dispatches ///////////////////////////////////// -void btOpenCLSoftBodySolver::copySoftBodyToVertexBuffer( const btSoftBody * const softBody, btVertexBufferDescriptor *vertexBuffer ) +void btSoftBodySolverOutputCLtoCPU::copySoftBodyToVertexBuffer( const btSoftBody * const softBody, btVertexBufferDescriptor *vertexBuffer ) { - // Currently only support CPU output buffers - // TODO: check for DX11 buffers. Take all offsets into the same DX11 buffer - // and use them together on a single kernel call if possible by setting up a - // per-cloth target buffer array for the copy kernel. + btSoftBodySolver *solver = softBody->getSoftBodySolver(); + btAssert( solver->getSolverType() == btSoftBodySolver::CL_SOLVER || solver->getSolverType() == btSoftBodySolver::CL_SIMD_SOLVER ); + btOpenCLSoftBodySolver *dxSolver = static_cast< btOpenCLSoftBodySolver * >( solver ); - btOpenCLAcceleratedSoftBodyInterface *currentCloth = findSoftBodyInterface( softBody ); + btOpenCLAcceleratedSoftBodyInterface* currentCloth = dxSolver->findSoftBodyInterface( softBody ); + btSoftBodyVertexDataOpenCL &vertexData( dxSolver->m_vertexData ); + const int firstVertex = currentCloth->getFirstVertex(); const int lastVertex = firstVertex + currentCloth->getNumVertices(); @@ -1180,8 +1521,8 @@ void btOpenCLSoftBodySolver::copySoftBodyToVertexBuffer( const btSoftBody * cons const btCPUVertexBufferDescriptor *cpuVertexBuffer = static_cast< btCPUVertexBufferDescriptor* >(vertexBuffer); float *basePointer = cpuVertexBuffer->getBasePointer(); - m_vertexData.m_clVertexPosition.copyFromGPU(); - m_vertexData.m_clVertexNormal.copyFromGPU(); + vertexData.m_clVertexPosition.copyFromGPU(); + vertexData.m_clVertexNormal.copyFromGPU(); if( vertexBuffer->hasVertexPositions() ) { @@ -1191,7 +1532,7 @@ void btOpenCLSoftBodySolver::copySoftBodyToVertexBuffer( const btSoftBody * cons for( int vertexIndex = firstVertex; vertexIndex < lastVertex; ++vertexIndex ) { - Vectormath::Aos::Point3 position = m_vertexData.getPosition(vertexIndex); + Vectormath::Aos::Point3 position = vertexData.getPosition(vertexIndex); *(vertexPointer + 0) = position.getX(); *(vertexPointer + 1) = position.getY(); *(vertexPointer + 2) = position.getZ(); @@ -1206,7 +1547,7 @@ void btOpenCLSoftBodySolver::copySoftBodyToVertexBuffer( const btSoftBody * cons for( int vertexIndex = firstVertex; vertexIndex < lastVertex; ++vertexIndex ) { - Vectormath::Aos::Vector3 normal = m_vertexData.getNormal(vertexIndex); + Vectormath::Aos::Vector3 normal = vertexData.getNormal(vertexIndex); *(normalPointer + 0) = normal.getX(); *(normalPointer + 1) = normal.getY(); *(normalPointer + 2) = normal.getZ(); @@ -1215,10 +1556,11 @@ void btOpenCLSoftBodySolver::copySoftBodyToVertexBuffer( const btSoftBody * cons } } -} // btCPUSoftBodySolver::outputToVertexBuffers +} // btSoftBodySolverOutputCLtoCPU::outputToVertexBuffers -cl_kernel btOpenCLSoftBodySolver::compileCLKernelFromString( const char* kernelSource, const char* kernelName ) + +cl_kernel CLFunctions::compileCLKernelFromString( const char* kernelSource, const char* kernelName, const char* additionalMacros ) { printf("compiling kernelName: %s ",kernelName); cl_kernel kernel; @@ -1229,19 +1571,45 @@ cl_kernel btOpenCLSoftBodySolver::compileCLKernelFromString( const char* kernelS // oclCHECKERROR(ciErrNum, CL_SUCCESS); // Build the program with 'mad' Optimization option + + #ifdef MAC char* flags = "-cl-mad-enable -DMAC -DGUID_ARG"; #else - const char* flags = "-DGUID_ARG="; + //const char* flags = "-DGUID_ARG= -fno-alias"; + const char* flags = "-DGUID_ARG= "; #endif - ciErrNum = clBuildProgram(m_cpProgram, 0, NULL, flags, NULL, NULL); + + char* compileFlags = new char[strlen(additionalMacros) + strlen(flags) + 5]; + sprintf(compileFlags, "%s %s", flags, additionalMacros); + ciErrNum = clBuildProgram(m_cpProgram, 0, NULL, compileFlags, NULL, NULL); if (ciErrNum != CL_SUCCESS) { - printf("Error in clBuildProgram, Line %u in file %s !!!\n\n", __LINE__, __FILE__); + size_t numDevices; + clGetProgramInfo( m_cpProgram, CL_PROGRAM_DEVICES, 0, 0, &numDevices ); + cl_device_id *devices = new cl_device_id[numDevices]; + clGetProgramInfo( m_cpProgram, CL_PROGRAM_DEVICES, numDevices, devices, &numDevices ); + for( int i = 0; i < 2; ++i ) + { + char *build_log; + size_t ret_val_size; + clGetProgramBuildInfo(m_cpProgram, devices[i], CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size); + build_log = new char[ret_val_size+1]; + clGetProgramBuildInfo(m_cpProgram, devices[i], CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL); + + // to be carefully, terminate with \0 + // there's no information in the reference whether the string is 0 terminated or not + build_log[ret_val_size] = '\0'; + + + printf("Error in clBuildProgram, Line %u in file %s, Log: \n%s\n !!!\n\n", __LINE__, __FILE__, build_log); + delete[] build_log; + } btAssert(0); exit(0); } + // Create the kernel kernel = clCreateKernel(m_cpProgram, kernelName, &ciErrNum); if (ciErrNum != CL_SUCCESS) @@ -1252,37 +1620,123 @@ cl_kernel btOpenCLSoftBodySolver::compileCLKernelFromString( const char* kernelS } printf("ready. \n"); + delete [] compileFlags; return kernel; } void btOpenCLSoftBodySolver::predictMotion( float timeStep ) { - // Fill the force arrays with current acceleration data etc - m_perClothWindVelocity.resize( m_softBodySet.size() ); - for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) + // Clear the collision shape array for the next frame + // Ensure that the DX11 ones are moved off the device so they will be updated correctly + m_clCollisionObjectDetails.changedOnCPU(); + m_clPerClothCollisionObjects.changedOnCPU(); + m_collisionObjectDetails.clear(); + { - btSoftBody *softBody = m_softBodySet[softBodyIndex]->getSoftBody(); - - m_perClothWindVelocity[softBodyIndex] = toVector3(softBody->getWindVelocity()); + BT_PROFILE("perClothWindVelocity"); + // Fill the force arrays with current acceleration data etc + m_perClothWindVelocity.resize( m_softBodySet.size() ); + for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) + { + btSoftBody *softBody = m_softBodySet[softBodyIndex]->getSoftBody(); + + m_perClothWindVelocity[softBodyIndex] = toVector3(softBody->getWindVelocity()); + } + } + { + BT_PROFILE("changedOnCPU"); + m_clPerClothWindVelocity.changedOnCPU(); } - m_clPerClothWindVelocity.changedOnCPU(); - // Apply forces that we know about to the cloths - applyForces( timeStep * getTimeScale() ); + { + BT_PROFILE("applyForces"); + // Apply forces that we know about to the cloths + applyForces( timeStep * getTimeScale() ); + } - // Itegrate motion for all soft bodies dealt with by the solver - integrate( timeStep * getTimeScale() ); + { + BT_PROFILE("integrate"); + // Itegrate motion for all soft bodies dealt with by the solver + integrate( timeStep * getTimeScale() ); + } + + { + BT_PROFILE("updateBounds"); + updateBounds(); + } // End prediction work for solvers } +static Vectormath::Aos::Transform3 toTransform3( const btTransform &transform ) +{ + Vectormath::Aos::Transform3 outTransform; + outTransform.setCol(0, toVector3(transform.getBasis().getColumn(0))); + outTransform.setCol(1, toVector3(transform.getBasis().getColumn(1))); + outTransform.setCol(2, toVector3(transform.getBasis().getColumn(2))); + outTransform.setCol(3, toVector3(transform.getOrigin())); + return outTransform; +} + +void btOpenCLAcceleratedSoftBodyInterface::updateBounds( const btVector3 &lowerBound, const btVector3 &upperBound ) +{ + float scalarMargin = this->getSoftBody()->getCollisionShape()->getMargin(); + btVector3 vectorMargin( scalarMargin, scalarMargin, scalarMargin ); + m_softBody->m_bounds[0] = lowerBound - vectorMargin; + m_softBody->m_bounds[1] = upperBound + vectorMargin; +} // btOpenCLSoftBodySolver::btDX11AcceleratedSoftBodyInterface::updateBounds + +void btOpenCLSoftBodySolver::processCollision( btSoftBody*, btSoftBody* ) +{ + +} + +// Add the collision object to the set to deal with for a particular soft body +void btOpenCLSoftBodySolver::processCollision( btSoftBody *softBody, btCollisionObject* collisionObject ) +{ + int softBodyIndex = findSoftBodyIndex( softBody ); + + if( softBodyIndex >= 0 ) + { + btCollisionShape *collisionShape = collisionObject->getCollisionShape(); + float friction = collisionObject->getFriction(); + int shapeType = collisionShape->getShapeType(); + if( shapeType == CAPSULE_SHAPE_PROXYTYPE ) + { + // Add to the list of expected collision objects + CollisionShapeDescription newCollisionShapeDescription; + newCollisionShapeDescription.softBodyIdentifier = softBodyIndex; + newCollisionShapeDescription.collisionShapeType = shapeType; + // TODO: May need to transpose this matrix either here or in HLSL + newCollisionShapeDescription.shapeTransform = toTransform3(collisionObject->getWorldTransform()); + btCapsuleShape *capsule = static_cast( collisionShape ); + newCollisionShapeDescription.radius = capsule->getRadius(); + newCollisionShapeDescription.halfHeight = capsule->getHalfHeight(); + newCollisionShapeDescription.margin = capsule->getMargin(); + newCollisionShapeDescription.upAxis = capsule->getUpAxis(); + newCollisionShapeDescription.friction = friction; + btRigidBody* body = static_cast< btRigidBody* >( collisionObject ); + newCollisionShapeDescription.linearVelocity = toVector3(body->getLinearVelocity()); + newCollisionShapeDescription.angularVelocity = toVector3(body->getAngularVelocity()); + m_collisionObjectDetails.push_back( newCollisionShapeDescription ); + + } else { + btAssert("Unsupported collision shape type\n"); + } + } else { + btAssert("Unknown soft body"); + } +} // btOpenCLSoftBodySolver::processCollision -btOpenCLAcceleratedSoftBodyInterface *btOpenCLSoftBodySolver::findSoftBodyInterface( const btSoftBody* const softBody ) + + + +btOpenCLAcceleratedSoftBodyInterface* btOpenCLSoftBodySolver::findSoftBodyInterface( const btSoftBody* const softBody ) { for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) { - btOpenCLAcceleratedSoftBodyInterface *softBodyInterface = m_softBodySet[softBodyIndex]; + btOpenCLAcceleratedSoftBodyInterface* softBodyInterface = m_softBodySet[softBodyIndex]; if( softBodyInterface->getSoftBody() == softBody ) return softBodyInterface; } @@ -1290,27 +1744,50 @@ btOpenCLAcceleratedSoftBodyInterface *btOpenCLSoftBodySolver::findSoftBodyInterf } +int btOpenCLSoftBodySolver::findSoftBodyIndex( const btSoftBody* const softBody ) +{ + for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex ) + { + btOpenCLAcceleratedSoftBodyInterface* softBodyInterface = m_softBodySet[softBodyIndex]; + if( softBodyInterface->getSoftBody() == softBody ) + return softBodyIndex; + } + return 1; +} + +bool btOpenCLSoftBodySolver::checkInitialized() +{ + if( !m_shadersInitialized ) + if( buildShaders() ) + m_shadersInitialized = true; + + return m_shadersInitialized; +} + bool btOpenCLSoftBodySolver::buildShaders() { + // Ensure current kernels are released first + releaseKernels(); + bool returnVal = true; if( m_shadersInitialized ) return true; - prepareLinksKernel = compileCLKernelFromString( PrepareLinksCLString, "PrepareLinksKernel" ); - updatePositionsFromVelocitiesKernel = compileCLKernelFromString( UpdatePositionsFromVelocitiesCLString, "UpdatePositionsFromVelocitiesKernel" ); - solvePositionsFromLinksKernel = compileCLKernelFromString( SolvePositionsCLString, "SolvePositionsFromLinksKernel" ); - updateVelocitiesFromPositionsWithVelocitiesKernel = compileCLKernelFromString( UpdateNodesCLString, "updateVelocitiesFromPositionsWithVelocitiesKernel" ); - updateVelocitiesFromPositionsWithoutVelocitiesKernel = compileCLKernelFromString( UpdatePositionsCLString, "updateVelocitiesFromPositionsWithoutVelocitiesKernel" ); - integrateKernel = compileCLKernelFromString( IntegrateCLString, "IntegrateKernel" ); - applyForcesKernel = compileCLKernelFromString( ApplyForcesCLString, "ApplyForcesKernel" ); + prepareLinksKernel = clFunctions.compileCLKernelFromString( PrepareLinksCLString, "PrepareLinksKernel" ); + updatePositionsFromVelocitiesKernel = clFunctions.compileCLKernelFromString( UpdatePositionsFromVelocitiesCLString, "UpdatePositionsFromVelocitiesKernel" ); + solvePositionsFromLinksKernel = clFunctions.compileCLKernelFromString( SolvePositionsCLString, "SolvePositionsFromLinksKernel" ); + updateVelocitiesFromPositionsWithVelocitiesKernel = clFunctions.compileCLKernelFromString( UpdateNodesCLString, "updateVelocitiesFromPositionsWithVelocitiesKernel" ); + updateVelocitiesFromPositionsWithoutVelocitiesKernel = clFunctions.compileCLKernelFromString( UpdatePositionsCLString, "updateVelocitiesFromPositionsWithoutVelocitiesKernel" ); + computeBoundsKernel = clFunctions.compileCLKernelFromString( ComputeBoundsCLString, "ComputeBoundsKernel" ); + solveCollisionsAndUpdateVelocitiesKernel = clFunctions.compileCLKernelFromString( SolveCollisionsAndUpdateVelocitiesCLString, "SolveCollisionsAndUpdateVelocitiesKernel" ); + integrateKernel = clFunctions.compileCLKernelFromString( IntegrateCLString, "IntegrateKernel" ); + applyForcesKernel = clFunctions.compileCLKernelFromString( ApplyForcesCLString, "ApplyForcesKernel" ); // TODO: Rename to UpdateSoftBodies - resetNormalsAndAreasKernel = compileCLKernelFromString( UpdateNormalsCLString, "ResetNormalsAndAreasKernel" ); - normalizeNormalsAndAreasKernel = compileCLKernelFromString( UpdateNormalsCLString, "NormalizeNormalsAndAreasKernel" ); - updateSoftBodiesKernel = compileCLKernelFromString( UpdateNormalsCLString, "UpdateSoftBodiesKernel" ); - //outputToVertexArrayWithNormalsKernel = compileCLKernelFromString( OutputToVertexArrayCLString, "OutputToVertexArrayWithNormalsKernel" ); - //outputToVertexArrayWithoutNormalsKernel = compileCLKernelFromString( OutputToVertexArrayCLString, "OutputToVertexArrayWithoutNormalsKernel" ); + resetNormalsAndAreasKernel = clFunctions.compileCLKernelFromString( UpdateNormalsCLString, "ResetNormalsAndAreasKernel" ); + normalizeNormalsAndAreasKernel = clFunctions.compileCLKernelFromString( UpdateNormalsCLString, "NormalizeNormalsAndAreasKernel" ); + updateSoftBodiesKernel = clFunctions.compileCLKernelFromString( UpdateNormalsCLString, "UpdateSoftBodiesKernel" ); if( returnVal ) diff --git a/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCL.h b/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCL.h index ad6dda867..507143261 100644 --- a/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCL.h +++ b/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/btSoftBodySolver_OpenCL.h @@ -25,12 +25,60 @@ subject to the following restrictions: #include "btSoftBodySolverVertexData_OpenCL.h" #include "btSoftBodySolverTriangleData_OpenCL.h" +class CLFunctions +{ +protected: + cl_command_queue m_cqCommandQue; + cl_context m_cxMainContext; + +public: + CLFunctions(cl_command_queue cqCommandQue, cl_context cxMainContext) : + m_cqCommandQue( cqCommandQue ), + m_cxMainContext( cxMainContext ) + { + } + + + /** + * Compile a compute shader kernel from a string and return the appropriate cl_kernel object. + */ + cl_kernel compileCLKernelFromString( const char* kernelSource, const char* kernelName, const char* additionalMacros = "" ); +}; /** - * SoftBody class to maintain information about a soft body instance - * within a solver. - * This data addresses the main solver arrays. + * Entry in the collision shape array. + * Specifies the shape type, the transform matrix and the necessary details of the collisionShape. */ +struct CollisionShapeDescription +{ + Vectormath::Aos::Transform3 shapeTransform; + Vectormath::Aos::Vector3 linearVelocity; + Vectormath::Aos::Vector3 angularVelocity; + + int softBodyIdentifier; + int collisionShapeType; + + // Both needed for capsule + float radius; + float halfHeight; + int upAxis; + + float margin; + float friction; + + CollisionShapeDescription() + { + collisionShapeType = 0; + margin = 0; + friction = 0; + } +}; + +/** + * SoftBody class to maintain information about a soft body instance + * within a solver. + * This data addresses the main solver arrays. + */ class btOpenCLAcceleratedSoftBodyInterface { protected: @@ -100,6 +148,11 @@ public: { return m_firstTriangle; } + + /** + * Update the bounds in the btSoftBody object + */ + void updateBounds( const btVector3 &lowerBound, const btVector3 &upperBound ); // TODO: All of these set functions will have to do checks and // update the world because restructuring of the arrays will be necessary @@ -108,7 +161,7 @@ public: { m_numVertices = numVertices; } - + void setNumTriangles( int numTriangles ) { m_numTriangles = numTriangles; @@ -172,20 +225,61 @@ public: }; + class btOpenCLSoftBodySolver : public btSoftBodySolver { -private: +public: + + + struct UIntVector3 + { + UIntVector3() + { + x = 0; + y = 0; + z = 0; + _padding = 0; + } + + UIntVector3( unsigned int x_, unsigned int y_, unsigned int z_ ) + { + x = x_; + y = y_; + z = z_; + _padding = 0; + } + + unsigned int x; + unsigned int y; + unsigned int z; + unsigned int _padding; + }; + + struct CollisionObjectIndices + { + CollisionObjectIndices( int f, int e ) + { + firstObject = f; + endObject = e; + } + + int firstObject; + int endObject; + }; btSoftBodyLinkDataOpenCL m_linkData; btSoftBodyVertexDataOpenCL m_vertexData; btSoftBodyTriangleDataOpenCL m_triangleData; +protected: + + CLFunctions clFunctions; + /** Variable to define whether we need to update solver constants on the next iteration */ bool m_updateSolverConstants; bool m_shadersInitialized; - /** * Cloths owned by this solver. * Only our cloths are in this array. @@ -224,6 +318,46 @@ private: btAlignedObjectArray< float > m_perClothMediumDensity; btOpenCLBuffer m_clPerClothMediumDensity; + /** + * Collision shape details: pair of index of first collision shape for the cloth and number of collision objects. + */ + btAlignedObjectArray< CollisionObjectIndices > m_perClothCollisionObjects; + btOpenCLBuffer m_clPerClothCollisionObjects; + + /** + * Collision shapes being passed across to the cloths in this solver. + */ + btAlignedObjectArray< CollisionShapeDescription > m_collisionObjectDetails; + btOpenCLBuffer< CollisionShapeDescription > m_clCollisionObjectDetails; + + /** + * Minimum bounds for each cloth. + * Updated by GPU and returned for use by broad phase. + * These are int vectors as a reminder that they store the int representation of a float, not a float. + * Bit 31 is inverted - is floats are stored with int-sortable values. + * This is really a uint4 array but thanks to a limitation of OpenCL atomics we are using uints. + */ + btAlignedObjectArray< UIntVector3 > m_perClothMinBounds; + btOpenCLBuffer< UIntVector3 > m_clPerClothMinBounds; + + /** + * Maximum bounds for each cloth. + * Updated by GPU and returned for use by broad phase. + * These are int vectors as a reminder that they store the int representation of a float, not a float. + * Bit 31 is inverted - is floats are stored with int-sortable values. + */ + btAlignedObjectArray< UIntVector3 > m_perClothMaxBounds; + btOpenCLBuffer< UIntVector3 > m_clPerClothMaxBounds; + + + /** + * Friction coefficient for each cloth + */ + btAlignedObjectArray< float > m_perClothFriction; + btOpenCLBuffer< float > m_clPerClothFriction; + + + cl_kernel prepareLinksKernel; cl_kernel solvePositionsFromLinksKernel; cl_kernel updateConstantsKernel; @@ -233,41 +367,37 @@ private: cl_kernel updateVelocitiesFromPositionsWithoutVelocitiesKernel; cl_kernel updateVelocitiesFromPositionsWithVelocitiesKernel; cl_kernel vSolveLinksKernel; + cl_kernel solveCollisionsAndUpdateVelocitiesKernel; cl_kernel resetNormalsAndAreasKernel; cl_kernel normalizeNormalsAndAreasKernel; + cl_kernel computeBoundsKernel; cl_kernel updateSoftBodiesKernel; - cl_kernel outputToVertexArrayWithNormalsKernel; - cl_kernel outputToVertexArrayWithoutNormalsKernel; cl_kernel outputToVertexArrayKernel; cl_kernel applyForcesKernel; - cl_kernel collideSphereKernel; - cl_kernel collideCylinderKernel; cl_command_queue m_cqCommandQue; cl_context m_cxMainContext; - + size_t m_defaultWorkGroupSize; - /** - * Compile a compute shader kernel from a string and return the appropriate cl_kernel object. - */ - cl_kernel compileCLKernelFromString( const char *shaderString, const char *shaderName ); - - bool buildShaders(); + virtual bool buildShaders(); void resetNormalsAndAreas( int numVertices ); void normalizeNormalsAndAreas( int numVertices ); void executeUpdateSoftBodies( int firstTriangle, int numTriangles ); + + void prepareCollisionConstraints(); Vectormath::Aos::Vector3 ProjectOnAxis( const Vectormath::Aos::Vector3 &v, const Vectormath::Aos::Vector3 &a ); void ApplyClampedForce( float solverdt, const Vectormath::Aos::Vector3 &force, const Vectormath::Aos::Vector3 &vertexVelocity, float inverseMass, Vectormath::Aos::Vector3 &vertexForce ); - btOpenCLAcceleratedSoftBodyInterface *findSoftBodyInterface( const btSoftBody* const softBody ); + + int findSoftBodyIndex( const btSoftBody* const softBody ); virtual void applyForces( float solverdt ); @@ -276,7 +406,7 @@ private: */ virtual void integrate( float solverdt ); - void updateConstants( float timeStep ); + virtual void updateConstants( float timeStep ); float computeTriangleArea( const Vectormath::Aos::Point3 &vertex0, @@ -292,15 +422,20 @@ private: void updatePositionsFromVelocities( float solverdt ); - void solveLinksForPosition( int startLink, int numLinks, float kst, float ti ); + virtual void solveLinksForPosition( int startLink, int numLinks, float kst, float ti ); void updateVelocitiesFromPositionsWithVelocities( float isolverdt ); void updateVelocitiesFromPositionsWithoutVelocities( float isolverdt ); + void computeBounds( ); + virtual void solveCollisionsAndUpdateVelocities( float isolverdt ); // End kernel dispatches ///////////////////////////////////// + void updateBounds(); + + void releaseKernels(); public: btOpenCLSoftBodySolver(cl_command_queue queue,cl_context ctx); @@ -308,7 +443,8 @@ public: virtual ~btOpenCLSoftBodySolver(); - + + btOpenCLAcceleratedSoftBodyInterface *findSoftBodyInterface( const btSoftBody* const softBody ); virtual btSoftBodyLinkData &getLinkData(); @@ -316,20 +452,27 @@ public: virtual btSoftBodyTriangleData &getTriangleData(); - + virtual SolverTypes getSolverType() const + { + return CL_SOLVER; + } virtual bool checkInitialized(); virtual void updateSoftBodies( ); - virtual void optimize( btAlignedObjectArray< btSoftBody * > &softBodies ); + virtual void optimize( btAlignedObjectArray< btSoftBody * > &softBodies , bool forceUpdate=false); + + virtual void copyBackToSoftBodies(); virtual void solveConstraints( float solverdt ); virtual void predictMotion( float solverdt ); - virtual void copySoftBodyToVertexBuffer( const btSoftBody *const softBody, btVertexBufferDescriptor *vertexBuffer ); + virtual void processCollision( btSoftBody *, btCollisionObject* ); + + virtual void processCollision( btSoftBody*, btSoftBody* ); virtual void setDefaultWorkgroupSize(size_t workGroupSize) { @@ -339,6 +482,27 @@ public: { return m_defaultWorkGroupSize; } + }; // btOpenCLSoftBodySolver + +/** + * Class to manage movement of data from a solver to a given target. + * This version is the CL to CPU version. + */ +class btSoftBodySolverOutputCLtoCPU : public btSoftBodySolverOutput +{ +protected: + +public: + btSoftBodySolverOutputCLtoCPU() + { + } + + /** Output current computed vertex data to the vertex buffers for all cloths in the solver. */ + virtual void copySoftBodyToVertexBuffer( const btSoftBody * const softBody, btVertexBufferDescriptor *vertexBuffer ); +}; + + + #endif // #ifndef BT_SOFT_BODY_SOLVER_OPENCL_H diff --git a/src/BulletSoftBody/btDefaultSoftBodySolver.cpp b/src/BulletSoftBody/btDefaultSoftBodySolver.cpp index 903d1efe5..c876ebf1f 100644 --- a/src/BulletSoftBody/btDefaultSoftBodySolver.cpp +++ b/src/BulletSoftBody/btDefaultSoftBodySolver.cpp @@ -34,10 +34,13 @@ btDefaultSoftBodySolver::~btDefaultSoftBodySolver() { } +// In this case the data is already in the soft bodies so there is no need for us to do anything +void btDefaultSoftBodySolver::copyBackToSoftBodies() +{ +} - -void btDefaultSoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * > &softBodies ) +void btDefaultSoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * > &softBodies , bool forceUpdate) { m_softBodySet.copyFromArray( softBodies ); } @@ -121,6 +124,17 @@ void btDefaultSoftBodySolver::copySoftBodyToVertexBuffer( const btSoftBody *cons } } // btDefaultSoftBodySolver::copySoftBodyToVertexBuffer +void btDefaultSoftBodySolver::processCollision( btSoftBody* softBody, btSoftBody* otherSoftBody) +{ + softBody->defaultCollisionHandler( otherSoftBody); +} + +// For the default solver just leave the soft body to do its collision processing +void btDefaultSoftBodySolver::processCollision( btSoftBody *softBody, btCollisionObject* collisionObject ) +{ + softBody->defaultCollisionHandler( collisionObject ); +} // btDefaultSoftBodySolver::processCollision + void btDefaultSoftBodySolver::predictMotion( float timeStep ) { diff --git a/src/BulletSoftBody/btDefaultSoftBodySolver.h b/src/BulletSoftBody/btDefaultSoftBodySolver.h index 93c1ff0d3..8e7db3daf 100644 --- a/src/BulletSoftBody/btDefaultSoftBodySolver.h +++ b/src/BulletSoftBody/btDefaultSoftBodySolver.h @@ -34,18 +34,30 @@ public: btDefaultSoftBodySolver(); virtual ~btDefaultSoftBodySolver(); + + virtual SolverTypes getSolverType() const + { + return DEFAULT_SOLVER; + } virtual bool checkInitialized(); virtual void updateSoftBodies( ); - virtual void optimize( btAlignedObjectArray< btSoftBody * > &softBodies ); + virtual void optimize( btAlignedObjectArray< btSoftBody * > &softBodies,bool forceUpdate=false ); + + virtual void copyBackToSoftBodies(); virtual void solveConstraints( float solverdt ); virtual void predictMotion( float solverdt ); virtual void copySoftBodyToVertexBuffer( const btSoftBody *const softBody, btVertexBufferDescriptor *vertexBuffer ); + + virtual void processCollision( btSoftBody *, btCollisionObject* ); + + virtual void processCollision( btSoftBody*, btSoftBody* ); + }; #endif // #ifndef BT_ACCELERATED_SOFT_BODY_CPU_SOLVER_H diff --git a/src/BulletSoftBody/btSoftBody.cpp b/src/BulletSoftBody/btSoftBody.cpp index 33be38cdd..a90acb99f 100644 --- a/src/BulletSoftBody/btSoftBody.cpp +++ b/src/BulletSoftBody/btSoftBody.cpp @@ -21,7 +21,7 @@ subject to the following restrictions: // btSoftBody::btSoftBody(btSoftBodyWorldInfo* worldInfo,int node_count, const btVector3* x, const btScalar* m) -:m_worldInfo(worldInfo) +:m_worldInfo(worldInfo),m_softBodySolver(0) { /* Init */ initDefaults(); @@ -2903,59 +2903,42 @@ btSoftBody::vsolver_t btSoftBody::getSolver(eVSolver::_ solver) // void btSoftBody::defaultCollisionHandler(btCollisionObject* pco) { -#if 0 - // If we have a solver, skip this work. - // It will have been done within the solver. - // TODO: In the case of the collision handler we need to ensure that we're - // updating the data structures correctly and in here we generate the - // collision lists to deal with in the solver - if( this->m_acceleratedSoftBody ) + + switch(m_cfg.collisions&fCollision::RVSmask) { - // We need to pass the collision data through to the solver collision engine - // Note that we only add the collision object here, we are not applying the collision or dealing with - // an impulse response. - m_acceleratedSoftBody->addCollisionObject(pco); - } else { -#endif - switch(m_cfg.collisions&fCollision::RVSmask) + case fCollision::SDF_RS: { - case fCollision::SDF_RS: - { - btSoftColliders::CollideSDF_RS docollide; - btRigidBody* prb1=btRigidBody::upcast(pco); - //btTransform wtr=prb1 ? prb1->getWorldTransform() : pco->getWorldTransform(); - btTransform wtr=pco->getWorldTransform(); + btSoftColliders::CollideSDF_RS docollide; + btRigidBody* prb1=btRigidBody::upcast(pco); + btTransform wtr=pco->getWorldTransform(); - const btTransform ctr=pco->getWorldTransform(); - const btScalar timemargin=(wtr.getOrigin()-ctr.getOrigin()).length(); - const btScalar basemargin=getCollisionShape()->getMargin(); - btVector3 mins; - btVector3 maxs; - ATTRIBUTE_ALIGNED16(btDbvtVolume) volume; - pco->getCollisionShape()->getAabb( pco->getWorldTransform(), - mins, - maxs); - volume=btDbvtVolume::FromMM(mins,maxs); - volume.Expand(btVector3(basemargin,basemargin,basemargin)); - docollide.psb = this; - docollide.m_colObj1 = pco; - docollide.m_rigidBody = prb1; + const btTransform ctr=pco->getWorldTransform(); + const btScalar timemargin=(wtr.getOrigin()-ctr.getOrigin()).length(); + const btScalar basemargin=getCollisionShape()->getMargin(); + btVector3 mins; + btVector3 maxs; + ATTRIBUTE_ALIGNED16(btDbvtVolume) volume; + pco->getCollisionShape()->getAabb( pco->getWorldTransform(), + mins, + maxs); + volume=btDbvtVolume::FromMM(mins,maxs); + volume.Expand(btVector3(basemargin,basemargin,basemargin)); + docollide.psb = this; + docollide.m_colObj1 = pco; + docollide.m_rigidBody = prb1; - docollide.dynmargin = basemargin+timemargin; - docollide.stamargin = basemargin; - m_ndbvt.collideTV(m_ndbvt.m_root,volume,docollide); - } - break; - case fCollision::CL_RS: - { - btSoftColliders::CollideCL_RS collider; - collider.Process(this,pco); - } - break; + docollide.dynmargin = basemargin+timemargin; + docollide.stamargin = basemargin; + m_ndbvt.collideTV(m_ndbvt.m_root,volume,docollide); } -#if 0 + break; + case fCollision::CL_RS: + { + btSoftColliders::CollideCL_RS collider; + collider.Process(this,pco); + } + break; } -#endif } // diff --git a/src/BulletSoftBody/btSoftBody.h b/src/BulletSoftBody/btSoftBody.h index f34d9f678..87247c3c1 100644 --- a/src/BulletSoftBody/btSoftBody.h +++ b/src/BulletSoftBody/btSoftBody.h @@ -37,7 +37,7 @@ subject to the following restrictions: class btBroadphaseInterface; class btDispatcher; - +class btSoftBodySolver; /* btSoftBodyWorldInfo */ struct btSoftBodyWorldInfo @@ -50,6 +50,17 @@ struct btSoftBodyWorldInfo btDispatcher* m_dispatcher; btVector3 m_gravity; btSparseSdf<3> m_sparsesdf; + + btSoftBodyWorldInfo() + :air_density((btScalar)1.2), + water_density(0), + water_offset(0), + water_normal(0,0,0), + m_broadphase(0), + m_dispatcher(0), + m_gravity(0,-10,0) + { + } }; @@ -60,6 +71,9 @@ class btSoftBody : public btCollisionObject public: btAlignedObjectArray m_collisionDisabledObjects; + // The solver object that handles this soft body + btSoftBodySolver *m_softBodySolver; + // // Enumerations // @@ -870,6 +884,31 @@ public: */ const btVector3& getWindVelocity(); + // + // Set the solver that handles this soft body + // Should not be allowed to get out of sync with reality + // Currently called internally on addition to the world + void setSoftBodySolver( btSoftBodySolver *softBodySolver ) + { + m_softBodySolver = softBodySolver; + } + + // + // Return the solver that handles this soft body + // + btSoftBodySolver *getSoftBodySolver() + { + return m_softBodySolver; + } + + // + // Return the solver that handles this soft body + // + btSoftBodySolver *getSoftBodySolver() const + { + return m_softBodySolver; + } + // // Cast diff --git a/src/BulletSoftBody/btSoftBodyInternals.h b/src/BulletSoftBody/btSoftBodyInternals.h index 04739e116..885571069 100644 --- a/src/BulletSoftBody/btSoftBodyInternals.h +++ b/src/BulletSoftBody/btSoftBodyInternals.h @@ -25,7 +25,7 @@ subject to the following restrictions: #include "BulletCollision/CollisionDispatch/btCollisionDispatcher.h" #include "BulletCollision/CollisionShapes/btConvexInternalShape.h" #include "BulletCollision/NarrowPhaseCollision/btGjkEpa2.h" - +#include //for memset // // btSymMatrix // @@ -172,8 +172,7 @@ public: template static inline void ZeroInitialize(T& value) { - static const T zerodummy; - value=zerodummy; + memset(&value,0,sizeof(T)); } // template diff --git a/src/BulletSoftBody/btSoftBodySolverVertexBuffer.h b/src/BulletSoftBody/btSoftBodySolverVertexBuffer.h index c4733d640..4c75bf216 100644 --- a/src/BulletSoftBody/btSoftBodySolverVertexBuffer.h +++ b/src/BulletSoftBody/btSoftBodySolverVertexBuffer.h @@ -1,165 +1,165 @@ -/* -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. -*/ - -#ifndef BT_SOFT_BODY_SOLVER_VERTEX_BUFFER_H -#define BT_SOFT_BODY_SOLVER_VERTEX_BUFFER_H - - -class btVertexBufferDescriptor -{ -public: - enum BufferTypes - { - CPU_BUFFER, - DX11_BUFFER, - OPENGL_BUFFER - }; - -protected: - - bool m_hasVertexPositions; - bool m_hasNormals; - - int m_vertexOffset; - int m_vertexStride; - - int m_normalOffset; - int m_normalStride; - -public: - btVertexBufferDescriptor() - { - m_hasVertexPositions = false; - m_hasNormals = false; - m_vertexOffset = 0; - m_vertexStride = 0; - m_normalOffset = 0; - m_normalStride = 0; - } - - virtual ~btVertexBufferDescriptor() - { - - } - - virtual bool hasVertexPositions() const - { - return m_hasVertexPositions; - } - - virtual bool hasNormals() const - { - return m_hasNormals; - } - - /** - * Return the type of the vertex buffer descriptor. - */ - virtual BufferTypes getBufferType() const = 0; - - /** - * Return the vertex offset in floats from the base pointer. - */ - virtual int getVertexOffset() const - { - return m_vertexOffset; - } - - /** - * Return the vertex stride in number of floats between vertices. - */ - virtual int getVertexStride() const - { - return m_vertexStride; - } - - /** - * Return the vertex offset in floats from the base pointer. - */ - virtual int getNormalOffset() const - { - return m_normalOffset; - } - - /** - * Return the vertex stride in number of floats between vertices. - */ - virtual int getNormalStride() const - { - return m_normalStride; - } -}; - - -class btCPUVertexBufferDescriptor : public btVertexBufferDescriptor -{ -protected: - float *m_basePointer; - -public: - /** - * vertexBasePointer is pointer to beginning of the buffer. - * vertexOffset is the offset in floats to the first vertex. - * vertexStride is the stride in floats between vertices. - */ - btCPUVertexBufferDescriptor( float *basePointer, int vertexOffset, int vertexStride ) - { - m_basePointer = basePointer; - m_vertexOffset = vertexOffset; - m_vertexStride = vertexStride; - m_hasVertexPositions = true; - } - - /** - * vertexBasePointer is pointer to beginning of the buffer. - * vertexOffset is the offset in floats to the first vertex. - * vertexStride is the stride in floats between vertices. - */ - btCPUVertexBufferDescriptor( float *basePointer, int vertexOffset, int vertexStride, int normalOffset, int normalStride ) - { - m_basePointer = basePointer; - - m_vertexOffset = vertexOffset; - m_vertexStride = vertexStride; - m_hasVertexPositions = true; - - m_normalOffset = normalOffset; - m_normalStride = normalStride; - m_hasNormals = true; - } - - virtual ~btCPUVertexBufferDescriptor() - { - - } - - /** - * Return the type of the vertex buffer descriptor. - */ - virtual BufferTypes getBufferType() const - { - return CPU_BUFFER; - } - - /** - * Return the base pointer in memory to the first vertex. - */ - virtual float *getBasePointer() const - { - return m_basePointer; - } -}; - -#endif // #ifndef BT_SOFT_BODY_SOLVER_VERTEX_BUFFER_H +/* +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. +*/ + +#ifndef BT_SOFT_BODY_SOLVER_VERTEX_BUFFER_H +#define BT_SOFT_BODY_SOLVER_VERTEX_BUFFER_H + + +class btVertexBufferDescriptor +{ +public: + enum BufferTypes + { + CPU_BUFFER, + DX11_BUFFER, + OPENGL_BUFFER + }; + +protected: + + bool m_hasVertexPositions; + bool m_hasNormals; + + int m_vertexOffset; + int m_vertexStride; + + int m_normalOffset; + int m_normalStride; + +public: + btVertexBufferDescriptor() + { + m_hasVertexPositions = false; + m_hasNormals = false; + m_vertexOffset = 0; + m_vertexStride = 0; + m_normalOffset = 0; + m_normalStride = 0; + } + + virtual ~btVertexBufferDescriptor() + { + + } + + virtual bool hasVertexPositions() const + { + return m_hasVertexPositions; + } + + virtual bool hasNormals() const + { + return m_hasNormals; + } + + /** + * Return the type of the vertex buffer descriptor. + */ + virtual BufferTypes getBufferType() const = 0; + + /** + * Return the vertex offset in floats from the base pointer. + */ + virtual int getVertexOffset() const + { + return m_vertexOffset; + } + + /** + * Return the vertex stride in number of floats between vertices. + */ + virtual int getVertexStride() const + { + return m_vertexStride; + } + + /** + * Return the vertex offset in floats from the base pointer. + */ + virtual int getNormalOffset() const + { + return m_normalOffset; + } + + /** + * Return the vertex stride in number of floats between vertices. + */ + virtual int getNormalStride() const + { + return m_normalStride; + } +}; + + +class btCPUVertexBufferDescriptor : public btVertexBufferDescriptor +{ +protected: + float *m_basePointer; + +public: + /** + * vertexBasePointer is pointer to beginning of the buffer. + * vertexOffset is the offset in floats to the first vertex. + * vertexStride is the stride in floats between vertices. + */ + btCPUVertexBufferDescriptor( float *basePointer, int vertexOffset, int vertexStride ) + { + m_basePointer = basePointer; + m_vertexOffset = vertexOffset; + m_vertexStride = vertexStride; + m_hasVertexPositions = true; + } + + /** + * vertexBasePointer is pointer to beginning of the buffer. + * vertexOffset is the offset in floats to the first vertex. + * vertexStride is the stride in floats between vertices. + */ + btCPUVertexBufferDescriptor( float *basePointer, int vertexOffset, int vertexStride, int normalOffset, int normalStride ) + { + m_basePointer = basePointer; + + m_vertexOffset = vertexOffset; + m_vertexStride = vertexStride; + m_hasVertexPositions = true; + + m_normalOffset = normalOffset; + m_normalStride = normalStride; + m_hasNormals = true; + } + + virtual ~btCPUVertexBufferDescriptor() + { + + } + + /** + * Return the type of the vertex buffer descriptor. + */ + virtual BufferTypes getBufferType() const + { + return CPU_BUFFER; + } + + /** + * Return the base pointer in memory to the first vertex. + */ + virtual float *getBasePointer() const + { + return m_basePointer; + } +}; + +#endif // #ifndef BT_SOFT_BODY_SOLVER_VERTEX_BUFFER_H diff --git a/src/BulletSoftBody/btSoftBodySolvers.h b/src/BulletSoftBody/btSoftBodySolvers.h index 79a55f706..01168b613 100644 --- a/src/BulletSoftBody/btSoftBodySolvers.h +++ b/src/BulletSoftBody/btSoftBodySolvers.h @@ -1,142 +1,154 @@ -/* -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. -*/ - -#ifndef BT_SOFT_BODY_SOLVERS_H -#define BT_SOFT_BODY_SOLVERS_H - -#include "BulletCollision/CollisionShapes/btTriangleIndexVertexArray.h" - - -class btSoftBodyTriangleData; -class btSoftBodyLinkData; -class btSoftBodyVertexData; -class btVertexBufferDescriptor; -class btCollisionObject; -class btSoftBody; - - -class btSoftBodySolver -{ - -protected: - int m_numberOfPositionIterations; - int m_numberOfVelocityIterations; - // Simulation timescale - float m_timeScale; - -public: - btSoftBodySolver() : - m_numberOfPositionIterations( 10 ), - m_timeScale( 1 ) - { - m_numberOfVelocityIterations = 0; - m_numberOfPositionIterations = 5; - } - - virtual ~btSoftBodySolver() - { - } - -#if 0 - /** Acceleration for all cloths in the solver. Can be used to efficiently apply gravity. */ - virtual void setPerClothAcceleration( int clothIdentifier, Vectormath::Aos::Vector3 acceleration ) = 0; - - /** A wind velocity applied normal to the cloth for all cloths in the solver. */ - virtual void setPerClothWindVelocity( int clothIdentifier, Vectormath::Aos::Vector3 windVelocity ) = 0; - - /** Set the density of the medium a given cloth is situated in. This could be air or possibly water. */ - virtual void setPerClothMediumDensity( int clothIdentifier, float mediumDensity ) = 0; - - /** A damping factor specific to each cloth applied for all cloths. */ - virtual void setPerClothDampingFactor( int clothIdentifier, float dampingFactor ) = 0; - - /** A damping factor specific to each cloth applied for all cloths. */ - virtual void setPerClothVelocityCorrectionCoefficient( int clothIdentifier, float velocityCorrectionCoefficient ) = 0; - - /** Lift parameter for wind action on cloth. */ - virtual void setPerClothLiftFactor( int clothIdentifier, float liftFactor ) = 0; - - /** Drag parameter for wind action on cloth. */ - virtual void setPerClothDragFactor( int clothIdentifier, float dragFactor ) = 0; - - /** - * Add a velocity to all soft bodies in the solver - useful for doing world-wide velocities such as a change due to gravity - * Only add a velocity to nodes with a non-zero inverse mass. - */ - virtual void addVelocity( Vectormath::Aos::Vector3 velocity ) = 0; -#endif - - - - /** Ensure that this solver is initialized. */ - virtual bool checkInitialized() = 0; - - /** Optimize soft bodies in this solver. */ - virtual void optimize( btAlignedObjectArray< btSoftBody * > &softBodies ) = 0; - - /** Predict motion of soft bodies into next timestep */ - virtual void predictMotion( float solverdt ) = 0; - - /** Solve constraints for a set of soft bodies */ - virtual void solveConstraints( float solverdt ) = 0; - - /** Perform necessary per-step updates of soft bodies such as recomputing normals and bounding boxes */ - virtual void updateSoftBodies() = 0; - - /** Output current computed vertex data to the vertex buffers for all cloths in the solver. */ - virtual void copySoftBodyToVertexBuffer( const btSoftBody * const softBody, btVertexBufferDescriptor *vertexBuffer ) = 0; - - - - /** Set the number of velocity constraint solver iterations this solver uses. */ - virtual void setNumberOfPositionIterations( int iterations ) - { - m_numberOfPositionIterations = iterations; - } - - /** Get the number of velocity constraint solver iterations this solver uses. */ - virtual int getNumberOfPositionIterations() - { - return m_numberOfPositionIterations; - } - - /** Set the number of velocity constraint solver iterations this solver uses. */ - virtual void setNumberOfVelocityIterations( int iterations ) - { - m_numberOfVelocityIterations = iterations; - } - - /** Get the number of velocity constraint solver iterations this solver uses. */ - virtual int getNumberOfVelocityIterations() - { - return m_numberOfVelocityIterations; - } - - /** Return the timescale that the simulation is using */ - float getTimeScale() - { - return m_timeScale; - } - -#if 0 - /** - * Add a collision object to be used by the indicated softbody. - */ - virtual void addCollisionObjectForSoftBody( int clothIdentifier, btCollisionObject *collisionObject ) = 0; -#endif -}; - - -#endif // #ifndef BT_SOFT_BODY_SOLVERS_H +/* +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. +*/ + +#ifndef BT_SOFT_BODY_SOLVERS_H +#define BT_SOFT_BODY_SOLVERS_H + +#include "BulletCollision/CollisionShapes/btTriangleIndexVertexArray.h" + + +class btSoftBodyTriangleData; +class btSoftBodyLinkData; +class btSoftBodyVertexData; +class btVertexBufferDescriptor; +class btCollisionObject; +class btSoftBody; + + +class btSoftBodySolver +{ +public: + enum SolverTypes + { + DEFAULT_SOLVER, + CPU_SOLVER, + CL_SOLVER, + CL_SIMD_SOLVER, + DX_SOLVER, + DX_SIMD_SOLVER + }; + + +protected: + int m_numberOfPositionIterations; + int m_numberOfVelocityIterations; + // Simulation timescale + float m_timeScale; + +public: + btSoftBodySolver() : + m_numberOfPositionIterations( 10 ), + m_timeScale( 1 ) + { + m_numberOfVelocityIterations = 0; + m_numberOfPositionIterations = 5; + } + + virtual ~btSoftBodySolver() + { + } + + /** + * Return the type of the solver. + */ + virtual SolverTypes getSolverType() const = 0; + + + /** Ensure that this solver is initialized. */ + virtual bool checkInitialized() = 0; + + /** Optimize soft bodies in this solver. */ + virtual void optimize( btAlignedObjectArray< btSoftBody * > &softBodies , bool forceUpdate=false) = 0; + + /** Copy necessary data back to the original soft body source objects. */ + virtual void copyBackToSoftBodies() = 0; + + /** Predict motion of soft bodies into next timestep */ + virtual void predictMotion( float solverdt ) = 0; + + /** Solve constraints for a set of soft bodies */ + virtual void solveConstraints( float solverdt ) = 0; + + /** Perform necessary per-step updates of soft bodies such as recomputing normals and bounding boxes */ + virtual void updateSoftBodies() = 0; + + /** Process a collision between one of the world's soft bodies and another collision object */ + virtual void processCollision( btSoftBody *, btCollisionObject* ) = 0; + + /** Process a collision between two soft bodies */ + virtual void processCollision( btSoftBody*, btSoftBody* ) = 0; + + /** Set the number of velocity constraint solver iterations this solver uses. */ + virtual void setNumberOfPositionIterations( int iterations ) + { + m_numberOfPositionIterations = iterations; + } + + /** Get the number of velocity constraint solver iterations this solver uses. */ + virtual int getNumberOfPositionIterations() + { + return m_numberOfPositionIterations; + } + + /** Set the number of velocity constraint solver iterations this solver uses. */ + virtual void setNumberOfVelocityIterations( int iterations ) + { + m_numberOfVelocityIterations = iterations; + } + + /** Get the number of velocity constraint solver iterations this solver uses. */ + virtual int getNumberOfVelocityIterations() + { + return m_numberOfVelocityIterations; + } + + /** Return the timescale that the simulation is using */ + float getTimeScale() + { + return m_timeScale; + } + +#if 0 + /** + * Add a collision object to be used by the indicated softbody. + */ + virtual void addCollisionObjectForSoftBody( int clothIdentifier, btCollisionObject *collisionObject ) = 0; +#endif +}; + +/** + * Class to manage movement of data from a solver to a given target. + * This version is abstract. Subclasses will have custom pairings for different combinations. + */ +class btSoftBodySolverOutput +{ +protected: + +public: + btSoftBodySolverOutput() + { + } + + virtual ~btSoftBodySolverOutput() + { + } + + + /** Output current computed vertex data to the vertex buffers for all cloths in the solver. */ + virtual void copySoftBodyToVertexBuffer( const btSoftBody * const softBody, btVertexBufferDescriptor *vertexBuffer ) = 0; +}; + + +#endif // #ifndef BT_SOFT_BODY_SOLVERS_H diff --git a/src/BulletSoftBody/btSoftRigidCollisionAlgorithm.cpp b/src/BulletSoftBody/btSoftRigidCollisionAlgorithm.cpp index 11ad9e7da..bc374c805 100644 --- a/src/BulletSoftBody/btSoftRigidCollisionAlgorithm.cpp +++ b/src/BulletSoftBody/btSoftRigidCollisionAlgorithm.cpp @@ -19,6 +19,8 @@ subject to the following restrictions: #include "BulletCollision/CollisionShapes/btBoxShape.h" #include "BulletCollision/CollisionDispatch/btCollisionObject.h" #include "btSoftBody.h" +#include "BulletSoftBody/btSoftBodySolvers.h" + ///TODO: include all the shapes that the softbody can collide with ///alternatively, implement special case collision algorithms (just like for rigid collision shapes) @@ -61,7 +63,7 @@ void btSoftRigidCollisionAlgorithm::processCollision (btCollisionObject* body0,b if (softBody->m_collisionDisabledObjects.findLinearSearch(rigidCollisionObject)==softBody->m_collisionDisabledObjects.size()) { - softBody->defaultCollisionHandler(rigidCollisionObject); + softBody->getSoftBodySolver()->processCollision(softBody, rigidCollisionObject); } diff --git a/src/BulletSoftBody/btSoftRigidDynamicsWorld.cpp b/src/BulletSoftBody/btSoftRigidDynamicsWorld.cpp index dba41436e..1b9b5e392 100644 --- a/src/BulletSoftBody/btSoftRigidDynamicsWorld.cpp +++ b/src/BulletSoftBody/btSoftRigidDynamicsWorld.cpp @@ -82,14 +82,15 @@ void btSoftRigidDynamicsWorld::predictUnconstraintMotion(btScalar timeStep) void btSoftRigidDynamicsWorld::internalSingleStepSimulation( btScalar timeStep ) { + + // Let the solver grab the soft bodies and if necessary optimize for it + m_softBodySolver->optimize( getSoftBodyArray() ); + if( !m_softBodySolver->checkInitialized() ) { btAssert( "Solver initialization failed\n" ); } - // Let the solver grab the soft bodies and if necessary optimize for it - m_softBodySolver->optimize( getSoftBodyArray() ); - btDiscreteDynamicsWorld::internalSingleStepSimulation( timeStep ); ///solve soft bodies constraints @@ -128,6 +129,10 @@ void btSoftRigidDynamicsWorld::addSoftBody(btSoftBody* body,short int collisionF { m_softBodies.push_back(body); + // Set the soft body solver that will deal with this body + // to be the world's solver + body->setSoftBodySolver( m_softBodySolver ); + btCollisionWorld::addCollisionObject(body, collisionFilterGroup, collisionFilterMask); diff --git a/src/BulletSoftBody/btSoftSoftCollisionAlgorithm.cpp b/src/BulletSoftBody/btSoftSoftCollisionAlgorithm.cpp index 85a727944..1b8cfa723 100644 --- a/src/BulletSoftBody/btSoftSoftCollisionAlgorithm.cpp +++ b/src/BulletSoftBody/btSoftSoftCollisionAlgorithm.cpp @@ -17,6 +17,7 @@ subject to the following restrictions: #include "BulletCollision/CollisionDispatch/btCollisionDispatcher.h" #include "BulletCollision/CollisionShapes/btBoxShape.h" #include "BulletCollision/CollisionDispatch/btCollisionObject.h" +#include "BulletSoftBody/btSoftBodySolvers.h" #include "btSoftBody.h" #define USE_PERSISTENT_CONTACTS 1 @@ -36,7 +37,7 @@ void btSoftSoftCollisionAlgorithm::processCollision (btCollisionObject* body0,bt { btSoftBody* soft0 = (btSoftBody*)body0; btSoftBody* soft1 = (btSoftBody*)body1; - soft0->defaultCollisionHandler(soft1); + soft0->getSoftBodySolver()->processCollision(soft0, soft1); } btScalar btSoftSoftCollisionAlgorithm::calculateTimeOfImpact(btCollisionObject* /*body0*/,btCollisionObject* /*body1*/,const btDispatcherInfo& /*dispatchInfo*/,btManifoldResult* /*resultOut*/)