fixes in ray-convex for GPU

Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp

@@ -234,8 +234,8 @@ GpuRaytraceScene::GpuRaytraceScene()
 	m_raytraceData = new GpuRaytraceInternalData;
 
 	m_raytraceData->m_texId = new GLuint;
-	m_raytraceData->textureWidth = 1024;//1024;
+	m_raytraceData->textureWidth = 256;//1024;//1024;
-	m_raytraceData->textureHeight = 1024;
+	m_raytraceData->textureHeight = 256;//1024;
 
 	//create new texture
 	glGenTextures(1, m_raytraceData->m_texId);
@@ -471,9 +471,9 @@ void GpuRaytraceScene::renderScene()
 
 			if (hit)
 			{
-				m_raytraceData->m_texels[(i)*3+0] = 255;
+				m_raytraceData->m_texels[(i)*3+0] = 128+128.f*hits[i].m_hitNormal.x;
-				m_raytraceData->m_texels[(i)*3+1] = 0;
+				m_raytraceData->m_texels[(i)*3+1] = 128+128.f*hits[i].m_hitNormal.y;
-				m_raytraceData->m_texels[(i)*3+2] = 0;
+				m_raytraceData->m_texels[(i)*3+2] = 128+128.f*hits[i].m_hitNormal.z;
 			} else
 			{
 				m_raytraceData->m_texels[(i)*3+0] = 0;

src/Bullet3OpenCL/Initialize/b3OpenCLUtils.cpp

@@ -610,7 +610,7 @@ cl_program b3OpenCLUtils_compileCLProgramFromString(cl_context clContext, cl_dev
 	int fileUpToDate = 0;
 	int binaryFileValid=0;
 	
-	if (clFileNameForCaching)
+	if (!disableBinaryCaching && clFileNameForCaching)
 	{
 		clGetDeviceInfo(device, CL_DEVICE_NAME, 256, &deviceName, NULL);
 		clGetDeviceInfo(device, CL_DRIVER_VERSION, 256, &driverVersion, NULL);
@@ -918,7 +918,7 @@ cl_program b3OpenCLUtils_compileCLProgramFromString(cl_context clContext, cl_dev
 		}
 
 
-		if( clFileNameForCaching )
+		if( !disableBinaryCaching && clFileNameForCaching )
 		{	//	write to binary
 
 			cl_uint numAssociatedDevices;

src/Bullet3OpenCL/Raycast/b3GpuRaycast.cpp

@@ -209,6 +209,8 @@ void b3GpuRaycast::castRays(const b3AlignedObjectArray<b3RayInfo>& rays,	b3Align
 		int numBodies,const struct b3RigidBodyCL* bodies, int numCollidables, const struct b3Collidable* collidables, const struct b3GpuNarrowPhaseInternalData* narrowphaseData)
 {
 	
+	//castRaysHost(rays,hitResults,numBodies,bodies,numCollidables,collidables,narrowphaseData);
+
 	B3_PROFILE("castRaysGPU");
 
 	b3OpenCLArray<b3RayInfo> gpuRays(m_data->m_context,m_data->m_q);

src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.cl

@@ -14,12 +14,12 @@ typedef struct
 
 typedef struct
 {
-		float m_hitFraction;
+	float m_hitFraction;
-		int	m_hitResult0;
+	int	m_hitResult0;
-		int	m_hitResult1;
+	int	m_hitResult1;
-		int	m_hitResult2;
+	int	m_hitResult2;
-		float4	m_hitPoint;
+	float4	m_hitPoint;
-		float4	m_hitNormal;
+	float4	m_hitNormal;
 } b3RayHit;
 
 typedef struct
@@ -82,20 +82,18 @@ typedef struct
 typedef float4 Quaternion;
 
 __inline
-Quaternion qtMul(Quaternion a, Quaternion b);
+	Quaternion qtMul(Quaternion a, Quaternion b);
 
 __inline
-Quaternion qtNormalize(Quaternion in);
+	Quaternion qtNormalize(Quaternion in);
-
-__inline
-float4 qtRotate(Quaternion q, float4 vec);
-
-__inline
-Quaternion qtInvert(Quaternion q);
 
 
 __inline
-float dot3F4(float4 a, float4 b)
+	Quaternion qtInvert(Quaternion q);
+
+
+__inline
+	float dot3F4(float4 a, float4 b)
 {
 	float4 a1 = (float4)(a.xyz,0.f);
 	float4 b1 = (float4)(b.xyz,0.f);
@@ -104,65 +102,56 @@ float dot3F4(float4 a, float4 b)
 
 
 __inline
-Quaternion qtMul(Quaternion a, Quaternion b)
+	Quaternion qtMul(Quaternion a, Quaternion b)
 {
 	Quaternion ans;
 	ans = cross( a, b );
 	ans += a.w*b+b.w*a;
-//	ans.w = a.w*b.w - (a.x*b.x+a.y*b.y+a.z*b.z);
+	//	ans.w = a.w*b.w - (a.x*b.x+a.y*b.y+a.z*b.z);
 	ans.w = a.w*b.w - dot3F4(a, b);
 	return ans;
 }
 
 __inline
-Quaternion qtNormalize(Quaternion in)
+	Quaternion qtNormalize(Quaternion in)
 {
 	return fast_normalize(in);
-//	in /= length( in );
+	//	in /= length( in );
-//	return in;
+	//	return in;
 }
 __inline
-float4 qtRotate(Quaternion q, float4 vec)
+	float4 qtRotate(Quaternion q, float4 vec)
 {
 	Quaternion qInv = qtInvert( q );
 	float4 vcpy = vec;
 	vcpy.w = 0.f;
-	float4 out = qtMul(qtMul(q,vcpy),qInv);
+	float4 out = qtMul(q,vcpy);
+	out = qtMul(out,qInv);
 	return out;
 }
 
 __inline
-Quaternion qtInvert(Quaternion q)
+	Quaternion qtInvert(Quaternion q)
 {
 	return (Quaternion)(-q.xyz, q.w);
 }
 
 __inline
-float4 qtInvRotate(const Quaternion q, float4 vec)
+	float4 qtInvRotate(const Quaternion q, float4 vec)
 {
 	return qtRotate( qtInvert( q ), vec );
 }
 
-__inline
+
-float4 transform(const float4* p, const float4* translation, const Quaternion* orientation)
-{
-	return qtRotate( *orientation, *p ) + (*translation);
-}
 
 void	trInverse(float4 translationIn, Quaternion orientationIn,
-		float4* translationOut, Quaternion* orientationOut)
+	float4* translationOut, Quaternion* orientationOut)
 {
 	*orientationOut = qtInvert(orientationIn);
 	*translationOut = qtRotate(*orientationOut, -translationIn);
 }
 
-void	trMul(float4 translationA, Quaternion orientationA,
+
-						float4 translationB, Quaternion orientationB,
-		float4* translationOut, Quaternion* orientationOut)
-{
-	*orientationOut = qtMul(orientationA,orientationB);
-	*translationOut = transform(&translationB,&translationA,&orientationA);
-}
 
 
 
@@ -171,10 +160,10 @@ bool rayConvex(float4 rayFromLocal, float4 rayToLocal, int numFaces, int faceOff
 {
 	rayFromLocal.w = 0.f;
 	rayToLocal.w = 0.f;
-  bool result = true;
+	bool result = true;
 
-	float exitFraction = *hitFraction;
+	float exitFraction = hitFraction[0];
-	float enterFraction = -0.1f;
+	float enterFraction = -0.3f;
 	float4 curHitNormal = (float4)(0,0,0,0);
 	for (int i=0;i<numFaces && result;i++)
 	{
@@ -211,12 +200,15 @@ bool rayConvex(float4 rayFromLocal, float4 rayToLocal, int numFaces, int faceOff
 			result = false;
 	}
 
-	result = result && (enterFraction < 0.f);
+	if (enterFraction < 0.f)
+	{
+		result = false;
+	}
 
 	if (result)
 	{	
-		*hitFraction = enterFraction;
+		hitFraction[0] = enterFraction;
-		*hitNormal = curHitNormal;
+		hitNormal[0] = curHitNormal;
 	}
 	return result;
 }
@@ -238,26 +230,26 @@ bool sphere_intersect(float4 spherePos,  float radius, float4 rayFrom, float4 ra
 
 	float D = B * B - A*C;
 
-  if (D > 0.0)
+	if (D > 0.0f)
-  {
+	{
-     float t = (-B - sqrt(D))/A;
+		float t = (-B - sqrt(D))/A;
 
-        if ( (t >= 0.0f) && (t < (*hitFraction)) )
+		if ( (t >= 0.0f) && (t < (*hitFraction)) )
-        {
+		{
-					*hitFraction = t;
+			*hitFraction = t;
-          return true;
+			return true;
-				}
+		}
 	}
 	return false;
 }
 
 float4 setInterpolate3(float4 from, float4 to, float t)
 {
-		float s = 1.0f - t;
+	float s = 1.0f - t;
-		float4 result;
+	float4 result;
-		result = s * from + t * to;
+	result = s * from + t * to;
-		result.w = 0.f;	
+	result.w = 0.f;	
-		return result;	
+	return result;	
 }
 
 __kernel void rayCastKernel(  
@@ -272,72 +264,74 @@ __kernel void rayCastKernel(
 {
 
 	int i = get_global_id(0);
-	if (i<numRays)
+	if (i>=numRays)
+		return;
+
+	hitResults[i].m_hitFraction = 1.f;
+
+	float4 rayFrom = rays[i].m_from;
+	float4 rayTo = rays[i].m_to;
+	float hitFraction = 1.f;
+	float4 hitPoint;
+	float4 hitNormal;
+	int hitBodyIndex= -1;
+
+	int cachedCollidableIndex = -1;
+	Collidable cachedCollidable;
+
+	for (int b=0;b<numBodies;b++)
 	{
-		hitResults[i].m_hitFraction = 1.f;
 
-		float4 rayFrom = rays[i].m_from;
+		Body body = bodies[b];
-		float4 rayTo = rays[i].m_to;
+		float4 pos = body.m_pos;
-		float hitFraction = 1.f;
+		float4 orn = body.m_quat;
-		float4 hitPoint;
+		if (cachedCollidableIndex != body.m_collidableIdx)
-		float4 hitNormal;
-		int hitBodyIndex= -1;
-		
-		int cachedCollidableIndex = -1;		
-		Collidable cachedCollidable;
-		
-		for (int b=0;b<numBodies;b++)
 		{
-					
+			cachedCollidableIndex = body.m_collidableIdx;
-				float4 pos = bodies[b].m_pos;
+			cachedCollidable = collidables[cachedCollidableIndex];
-				float4 orn = bodies[b].m_quat;
-				if (cachedCollidableIndex !=bodies[b].m_collidableIdx)
-				{
-						cachedCollidableIndex = bodies[b].m_collidableIdx;
-						cachedCollidable = collidables[cachedCollidableIndex];
-				}
-				
-				if (cachedCollidable.m_shapeType == SHAPE_SPHERE)
-				{
-					float radius = cachedCollidable.m_radius;
-		
-					if (sphere_intersect(pos,  radius, rayFrom, rayTo, &hitFraction))
-					{
-						hitBodyIndex = b;
-						hitPoint = setInterpolate3(rayFrom, rayTo,hitFraction);
-						hitNormal = (float4) (hitPoint-bodies[b].m_pos);
-					}
-				}
-				
-				if (cachedCollidable.m_shapeType == SHAPE_CONVEX_HULL)
-				{
-				
-					float4 invPos = (float4)(0,0,0,0);
-					float4 invOrn = (float4)(0,0,0,0);
-					float4 rayFromLocal = (float4)(0,0,0,0);
-					float4 rayToLocal = (float4)(0,0,0,0);
-					
-					trInverse(pos,orn, &invPos, &invOrn);
-					rayFromLocal = transform(&rayFrom, &invPos, &invOrn);
-					rayToLocal = transform(&rayTo, &invPos, &invOrn);
-					
-					int numFaces = convexShapes[cachedCollidable.m_shapeIndex].m_numFaces;
-					int faceOffset = convexShapes[cachedCollidable.m_shapeIndex].m_faceOffset;
-					
-					if (rayConvex(rayFromLocal, rayToLocal, numFaces, faceOffset,faces, &hitFraction, &hitNormal))
-					{
-						hitBodyIndex = b;
-					}
-				}
-			
 		}
-		
+		if (cachedCollidable.m_shapeType == SHAPE_CONVEX_HULL)
-		if (hitBodyIndex>=0)
 		{
-			hitResults[i].m_hitFraction = hitFraction;
+
-			hitResults[i].m_hitPoint = hitPoint;
+			float4 invPos = (float4)(0,0,0,0);
-			hitResults[i].m_hitNormal = normalize(hitNormal);
+			float4 invOrn = (float4)(0,0,0,0);
-			hitResults[i].m_hitResult0 = hitBodyIndex;
+			float4 rayFromLocal = (float4)(0,0,0,0);
+			float4 rayToLocal = (float4)(0,0,0,0);
+			invOrn = qtInvert(orn);
+			invPos = qtRotate(invOrn, -pos);
+			rayFromLocal = qtRotate( invOrn, rayFrom ) + invPos;
+			rayToLocal = qtRotate( invOrn, rayTo) + invPos;
+			rayFromLocal.w = 0.f;
+			rayToLocal.w = 0.f;
+			int numFaces = convexShapes[cachedCollidable.m_shapeIndex].m_numFaces;
+			int faceOffset = convexShapes[cachedCollidable.m_shapeIndex].m_faceOffset;
+			if (numFaces)
+			{
+				if (rayConvex(rayFromLocal, rayToLocal, numFaces, faceOffset,faces, &hitFraction, &hitNormal))
+				{
+					hitBodyIndex = b;
+				}
+			}
+		}
+		if (cachedCollidable.m_shapeType == SHAPE_SPHERE)
+		{
+			float radius = cachedCollidable.m_radius;
+		
+			if (sphere_intersect(pos,  radius, rayFrom, rayTo, &hitFraction))
+			{
+				hitBodyIndex = b;
+				hitPoint = setInterpolate3(rayFrom, rayTo,hitFraction);
+				hitNormal = (float4) (hitPoint-bodies[b].m_pos);
+			}
 		}
 	}
+
+	if (hitBodyIndex>=0)
+	{
+		hitResults[i].m_hitFraction = hitFraction;
+		hitResults[i].m_hitPoint = hitPoint;
+		hitResults[i].m_hitNormal = normalize(hitNormal);
+		hitResults[i].m_hitResult0 = hitBodyIndex;
+	}
+
 }

src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h

@@ -176,7 +176,7 @@ static const char* rayCastKernelCL= \
 "  bool result = true;\n"
 "  \n"
 "	float exitFraction = *hitFraction;\n"
-"	float enterFraction = -0.1f;\n"
+"	float enterFraction = -0.3f;\n"
 "	float4 curHitNormal = (float4)(0,0,0,0);\n"
 "	for (int i=0;i<numFaces && result;i++)\n"
 "	{\n"
@@ -285,19 +285,17 @@ static const char* rayCastKernelCL= \
 "		float4 hitNormal;\n"
 "		int hitBodyIndex= -1;\n"
 "		\n"
-"		int cachedCollidableIndex = -1;		\n"
+"		\n"
-"		Collidable cachedCollidable;\n"
+"		\n"
 "		\n"
 "		for (int b=0;b<numBodies;b++)\n"
 "		{\n"
 "					\n"
 "				float4 pos = bodies[b].m_pos;\n"
 "				float4 orn = bodies[b].m_quat;\n"
-"				if (cachedCollidableIndex !=bodies[b].m_collidableIdx)\n"
+"		\n"
-"				{\n"
+"				int cachedCollidableIndex = bodies[b].m_collidableIdx;		\n"
-"						cachedCollidableIndex = bodies[b].m_collidableIdx;\n"
+"				Collidable cachedCollidable = collidables[cachedCollidableIndex];\n"
-"						cachedCollidable = collidables[cachedCollidableIndex];\n"
-"				}\n"
 "				\n"
 "				if (cachedCollidable.m_shapeType == SHAPE_SPHERE)\n"
 "				{\n"