Static objects support added for appGpu2dDemo

4 files changed, 43 insertions, 11 deletions

diff --git a/Demos/Gpu2dDemo/btGpuDemo2dSharedCode.h b/Demos/Gpu2dDemo/btGpuDemo2dSharedCode.h
index 0d0e44ca7..bcfece161 100644
--- a/Demos/Gpu2dDemo/btGpuDemo2dSharedCode.h
+++ b/Demos/Gpu2dDemo/btGpuDemo2dSharedCode.h
@@ -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));

diff --git a/Demos/Gpu2dDemo/btGpuDemo2dSharedDefs.h b/Demos/Gpu2dDemo/btGpuDemo2dSharedDefs.h
index cc72602e5..ff6e15d73 100644
--- a/Demos/Gpu2dDemo/btGpuDemo2dSharedDefs.h
+++ b/Demos/Gpu2dDemo/btGpuDemo2dSharedDefs.h
@@ -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"

 

diff --git a/Demos/Gpu2dDemo/btGpuDemoDynamicsWorld.cpp b/Demos/Gpu2dDemo/btGpuDemoDynamicsWorld.cpp
index f0b7bf8da..555ea247a 100644
--- a/Demos/Gpu2dDemo/btGpuDemoDynamicsWorld.cpp
+++ b/Demos/Gpu2dDemo/btGpuDemoDynamicsWorld.cpp
@@ -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;

 			}

diff --git a/Demos/Gpu2dDemo/btGpuDemoDynamicsWorld.h b/Demos/Gpu2dDemo/btGpuDemoDynamicsWorld.h
index a6c6fc9db..dfd4d98e6 100644
--- a/Demos/Gpu2dDemo/btGpuDemoDynamicsWorld.h
+++ b/Demos/Gpu2dDemo/btGpuDemoDynamicsWorld.h
@@ -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);