Bart/bullet3
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

Static objects support added for appGpu2dDemo

Browse Source
This commit is contained in:
rponom
2009-05-13 00:43:42 +00:00
parent 22fe8935c3
commit 4616b62686
4 changed files with 43 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files

26
Demos/Gpu2dDemo/btGpuDemo2dSharedCode.h
View File

@@ -202,6 +202,7 @@ BT_GPU___global__ void collisionWithWallBoxD(float4 *pos,
float *angVel,
char* shapes,
int2* shapeIds,
float* invMass,
btCudaPartProps pProp,
btCudaBoxProps gProp,
int nParticles,
@@ -214,8 +215,13 @@ BT_GPU___global__ void collisionWithWallBoxD(float4 *pos,
float3 impulse;
if((idx > 0) && idx < nParticles)
if((idx > 0) && (idx < nParticles))
{
float inv_mass = invMass[idx];
if(inv_mass <= 0.f)
{
return;
}
aPos=BT_GPU_make_float34(BT_GPU_FETCH4(pos,idx));
aRot=rotation[idx];
float4* shape = (float4*)(shapes + shapeIds[idx].x);
@@ -309,6 +315,7 @@ BT_GPU___device__ void collisionResolutionBox( int constrId,
float *angularVel,
float *lambdaDtBox,
float4* contact,
float* invMass,
btCudaPartProps pProp,
float dt)
{
@@ -366,12 +373,12 @@ BT_GPU___device__ void collisionResolutionBox( int constrId,
}
}
#endif //USE_FRICTION
if(aId)
if(aId && (invMass[aId] > 0.f))
{
aVel+= impulse;
aAngVel+= BT_GPU_cross(contactPoint, impulse).z;
}
if(bId)
if(bId && (invMass[bId] > 0.f))
{
bVel-= impulse;
bAngVel-= BT_GPU_cross(contactPoint+aPos-bPos, impulse).z;
@@ -394,6 +401,7 @@ BT_GPU___global__ void collisionBatchResolutionBoxD(int2 *constraints,
float *angularVel,
float *lambdaDtBox,
float4* contact,
float* invMass,
btCudaPartProps pProp,
int iBatch,
float dt)
@@ -403,7 +411,7 @@ BT_GPU___global__ void collisionBatchResolutionBoxD(int2 *constraints,
{
int idx = batch[k_idx];
collisionResolutionBox( idx, constraints, pos, vel, rotation, angularVel, lambdaDtBox,
contact, pProp, dt);
contact, invMass, pProp, dt);
}
}
@@ -448,7 +456,7 @@ void BT_GPU_PREF(setConstraintData(void* constraints,int numConstraints,int numO
BT_GPU_CHECK_ERROR("setConstraintDataD kernel execution failed");
}
void BT_GPU_PREF(collisionWithWallBox(void* pos,void* vel,float *rotation,float *angVel,char* shapes,void* shapeIds,btCudaPartProps pProp, btCudaBoxProps gProp,int numObjs,float dt))
void BT_GPU_PREF(collisionWithWallBox(void* pos,void* vel,float *rotation,float *angVel,char* shapes,void* shapeIds,void* invMass,btCudaPartProps pProp, btCudaBoxProps gProp,int numObjs,float dt))
{
if(!numObjs)
{
@@ -457,19 +465,20 @@ void BT_GPU_PREF(collisionWithWallBox(void* pos,void* vel,float *rotation,float
float4* pPos = (float4*)pos;
float4* pVel = (float4*)vel;
int2* pShapeIds = (int2*)shapeIds;
float* pInvMass = (float*)invMass;
BT_GPU_SAFE_CALL(BT_GPU_BindTexture(0, posTex, pPos, numObjs * sizeof(float4)));
int numThreads, numBlocks;
BT_GPU_PREF(computeGridSize)(numObjs, 256, numBlocks, numThreads);
// execute the kernel
BT_GPU_EXECKERNEL(numBlocks, numThreads, collisionWithWallBoxD, (pPos,pVel,rotation,angVel,shapes, pShapeIds,pProp,gProp,numObjs,dt));
BT_GPU_EXECKERNEL(numBlocks, numThreads, collisionWithWallBoxD, (pPos,pVel,rotation,angVel,shapes, pShapeIds,pInvMass,pProp,gProp,numObjs,dt));
BT_GPU_SAFE_CALL(BT_GPU_UnbindTexture(posTex));
// check if kernel invocation generated an error
BT_GPU_CHECK_ERROR("collisionWithWallBoxD kernel execution failed");
}
void BT_GPU_PREF(collisionBatchResolutionBox(void* constraints,int *batch,int numConstraints,int numObjs,void *pos,void *vel,float *rotation,float *angularVel,float *lambdaDtBox,void* contact,btCudaPartProps pProp,int iBatch,float dt))
void BT_GPU_PREF(collisionBatchResolutionBox(void* constraints,int *batch,int numConstraints,int numObjs,void *pos,void *vel,float *rotation,float *angularVel,float *lambdaDtBox,void* contact,void* invMass,btCudaPartProps pProp,int iBatch,float dt))
{
if(!numConstraints)
{
@@ -479,11 +488,12 @@ void BT_GPU_PREF(collisionBatchResolutionBox(void* constraints,int *batch,int nu
float4* pPos = (float4*)pos;
float4* pVel = (float4*)vel;
float4* pCont = (float4*)contact;
float* pInvMass = (float*)invMass;
int numThreads, numBlocks;
BT_GPU_PREF(computeGridSize)(numConstraints, 128, numBlocks, numThreads);
BT_GPU_SAFE_CALL(BT_GPU_BindTexture(0, posTex, pPos, numObjs * sizeof(float4)));
// execute the kernel
BT_GPU_EXECKERNEL(numBlocks, numThreads, collisionBatchResolutionBoxD, (pConstr,batch,numConstraints,pPos,pVel,rotation,angularVel,lambdaDtBox,pCont,pProp,iBatch,dt));
BT_GPU_EXECKERNEL(numBlocks, numThreads, collisionBatchResolutionBoxD, (pConstr,batch,numConstraints,pPos,pVel,rotation,angularVel,lambdaDtBox,pCont,pInvMass,pProp,iBatch,dt));
// check if kernel invocation generated an error
BT_GPU_CHECK_ERROR("collisionBatchResolutionBox2D kernel execution failed");
BT_GPU_SAFE_CALL(BT_GPU_UnbindTexture(posTex));

4
Demos/Gpu2dDemo/btGpuDemo2dSharedDefs.h
View File

@@ -25,8 +25,8 @@ extern "C"
void BT_GPU_PREF(clearAccumulationOfLambdaDt(float* lambdaDtBox, int numConstraints, int numContPoints));
void BT_GPU_PREF(setConstraintData(void* constraints,int numConstraints,int numObjs,void* pos,float *rotation,char* shapes,void* shapeIds,btCudaPartProps pProp,void* oContact));
void BT_GPU_PREF(collisionWithWallBox(void* pos,void* vel,float *rotation,float *angVel,char* shapes,void* shapeIds,btCudaPartProps pProp, btCudaBoxProps gProp,int numObjs,float dt));
void BT_GPU_PREF(collisionBatchResolutionBox(void* constraints,int *batch,int numConstraints,int numObjs,void *pos,void *vel,float *rotation,float *angularVel,float *lambdaDtBox,void* contact,btCudaPartProps pProp,int iBatch,float dt));
void BT_GPU_PREF(collisionWithWallBox(void* pos,void* vel,float *rotation,float *angVel,char* shapes,void* shapeIds,void* invMass,btCudaPartProps pProp,btCudaBoxProps gProp,int numObjs,float dt));
void BT_GPU_PREF(collisionBatchResolutionBox(void* constraints,int *batch,int numConstraints,int numObjs,void *pos,void *vel,float *rotation,float *angularVel,float *lambdaDtBox,void* contact,void* invMass,btCudaPartProps pProp,int iBatch,float dt));
} // extern "C"

12
Demos/Gpu2dDemo/btGpuDemoDynamicsWorld.cpp
View File

@@ -204,6 +204,10 @@ void btGpuDemoDynamicsWorld::grabData()
m_hVel[i+1] = *((float4*)&v);
v = rb->getAngularVelocity();
m_hAngVel[i+1] = v[2];
if(m_copyMassDataToGPU)
{
m_hInvMass[i+1] = rb->getInvMass();
}
}
if(m_useBulletNarrowphase)
{
@@ -328,6 +332,11 @@ void btGpuDemoDynamicsWorld::copyDataToGPU()
btCuda_copyArrayToDevice(m_dShapeIds, m_hShapeIds, (m_numObj + 1) * sizeof(int2));
m_copyShapeDataToGPU = false;
}
if(m_copyMassDataToGPU)
{
btCuda_copyArrayToDevice(m_dInvMass, m_hInvMass, (m_numObj + 1) * sizeof(float));
m_copyMassDataToGPU = false;
}
#endif //BT_USE_CUDA
} // btGpuDemoDynamicsWorld::copyDataToGPU()
@@ -465,7 +474,7 @@ void btGpuDemoDynamicsWorld::solveConstraintsCPU2(btContactSolverInfo& solverInf
for(int i=0;i<nIter;i++){
btGpu_collisionWithWallBox(m_hPos, m_hVel, m_hRot, m_hAngVel,m_hShapeBuffer, m_hShapeIds,
partProps, boxProps, m_numObj + 1, timeStep);
m_hInvMass, partProps, boxProps, m_numObj + 1, timeStep);
int* pBatchIds = m_hBatchIds;
for(int iBatch=0;iBatch < m_maxBatches;iBatch++)
{
@@ -479,6 +488,7 @@ void btGpuDemoDynamicsWorld::solveConstraintsCPU2(btContactSolverInfo& solverInf
m_hRot, m_hAngVel,
m_hLambdaDtBox,
m_hContact,
m_hInvMass,
partProps, iBatch, timeStep);
pBatchIds += numContConstraints;
}

12
Demos/Gpu2dDemo/btGpuDemoDynamicsWorld.h
View File

@@ -73,6 +73,10 @@ protected:
float4* m_hVel;
float* m_hAngVel;
float* m_hInvMass;
float* m_dInvMass;
bool m_copyMassDataToGPU;
#ifdef BT_USE_CUDA
float4* m_dPos;
float* m_dRot;
@@ -166,6 +170,8 @@ public:
m_hRot = new float[m_maxObjs];
m_hAngVel = new float[m_maxObjs];
m_hInvMass = new float[m_maxObjs];
m_maxVtxPerObj = 8;
#ifdef BT_USE_CUDA
@@ -178,6 +184,8 @@ public:
btCuda_allocateArray((void**)&m_dpVel, sizeof(float4) * m_maxObjs);
btCuda_allocateArray((void**)&m_dpAngVel, sizeof(float) * m_maxObjs);
btCuda_allocateArray((void**)&m_dInvMass, sizeof(float) * m_maxObjs);
btCuda_allocateArray((void**)&m_dIds, sizeof(int2) * sz);
btCuda_allocateArray((void**)&m_dBatchIds, sizeof(int) * sz);
@@ -205,6 +213,8 @@ public:
initShapeBuffer(m_maxObjs * CUDA_DEMO_DYNAMICS_WORLD_MAX_SPHERES_PER_OBJ * sizeof(float) * 4);
m_copyMassDataToGPU = true;
}
virtual ~btGpuDemoDynamicsWorld()
{
@@ -216,6 +226,7 @@ public:
delete [] m_hRot;
delete [] m_hVel;
delete [] m_hAngVel;
delete [] m_hInvMass;
#ifdef BT_USE_CUDA
btCuda_freeArray(m_dPos);
btCuda_freeArray(m_dRot);
@@ -225,6 +236,7 @@ public:
btCuda_freeArray(m_dpRot);
btCuda_freeArray(m_dpVel);
btCuda_freeArray(m_dpAngVel);
btCuda_freeArray(m_dInvMass);
btCuda_freeArray(m_dIds);
btCuda_freeArray(m_dBatchIds);