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);