From e39af464036546c88e5daabfd7dfa7031adb95a3 Mon Sep 17 00:00:00 2001 From: erwincoumans Date: Thu, 20 Jun 2013 16:41:11 -0700 Subject: [PATCH] fix ray-convex header file --- Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp | 4 +- .../Raycast/kernels/rayCastKernels.h | 238 +++++++++--------- 2 files changed, 119 insertions(+), 123 deletions(-) diff --git a/Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp b/Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp index b0989a2e8..699e1ffcf 100644 --- a/Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp +++ b/Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp @@ -234,8 +234,8 @@ GpuRaytraceScene::GpuRaytraceScene() m_raytraceData = new GpuRaytraceInternalData; m_raytraceData->m_texId = new GLuint; - m_raytraceData->textureWidth = 256;//1024;//1024; - m_raytraceData->textureHeight = 256;//1024; + m_raytraceData->textureWidth = 512;//1024;//1024; + m_raytraceData->textureHeight = 512;//1024; //create new texture glGenTextures(1, m_raytraceData->m_texId); diff --git a/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h b/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h index 02d8a43a4..f61793f01 100644 --- a/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h +++ b/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h @@ -16,12 +16,12 @@ static const char* rayCastKernelCL= \ "\n" "typedef struct\n" "{\n" -" float m_hitFraction;\n" -" int m_hitResult0;\n" -" int m_hitResult1;\n" -" int m_hitResult2;\n" -" float4 m_hitPoint;\n" -" float4 m_hitNormal;\n" +" float m_hitFraction;\n" +" int m_hitResult0;\n" +" int m_hitResult1;\n" +" int m_hitResult2;\n" +" float4 m_hitPoint;\n" +" float4 m_hitNormal;\n" "} b3RayHit;\n" "\n" "typedef struct\n" @@ -55,12 +55,12 @@ static const char* rayCastKernelCL= \ " float4 m_extents;\n" " float4 mC;\n" " float4 mE;\n" -" \n" +"\n" " float m_radius;\n" " int m_faceOffset;\n" " int m_numFaces;\n" " int m_numVertices;\n" -" \n" +"\n" " int m_vertexOffset;\n" " int m_uniqueEdgesOffset;\n" " int m_numUniqueEdges;\n" @@ -84,20 +84,18 @@ static const char* rayCastKernelCL= \ "typedef float4 Quaternion;\n" "\n" "__inline\n" -"Quaternion qtMul(Quaternion a, Quaternion b);\n" +" Quaternion qtMul(Quaternion a, Quaternion b);\n" "\n" "__inline\n" -"Quaternion qtNormalize(Quaternion in);\n" -"\n" -"__inline\n" -"float4 qtRotate(Quaternion q, float4 vec);\n" -"\n" -"__inline\n" -"Quaternion qtInvert(Quaternion q);\n" +" Quaternion qtNormalize(Quaternion in);\n" "\n" "\n" "__inline\n" -"float dot3F4(float4 a, float4 b)\n" +" Quaternion qtInvert(Quaternion q);\n" +"\n" +"\n" +"__inline\n" +" float dot3F4(float4 a, float4 b)\n" "{\n" " float4 a1 = (float4)(a.xyz,0.f);\n" " float4 b1 = (float4)(b.xyz,0.f);\n" @@ -106,65 +104,56 @@ static const char* rayCastKernelCL= \ "\n" "\n" "__inline\n" -"Quaternion qtMul(Quaternion a, Quaternion b)\n" +" Quaternion qtMul(Quaternion a, Quaternion b)\n" "{\n" " Quaternion ans;\n" " ans = cross( a, b );\n" " ans += a.w*b+b.w*a;\n" -"// ans.w = a.w*b.w - (a.x*b.x+a.y*b.y+a.z*b.z);\n" +" // ans.w = a.w*b.w - (a.x*b.x+a.y*b.y+a.z*b.z);\n" " ans.w = a.w*b.w - dot3F4(a, b);\n" " return ans;\n" "}\n" "\n" "__inline\n" -"Quaternion qtNormalize(Quaternion in)\n" +" Quaternion qtNormalize(Quaternion in)\n" "{\n" " return fast_normalize(in);\n" -"// in /= length( in );\n" -"// return in;\n" +" // in /= length( in );\n" +" // return in;\n" "}\n" "__inline\n" -"float4 qtRotate(Quaternion q, float4 vec)\n" +" float4 qtRotate(Quaternion q, float4 vec)\n" "{\n" " Quaternion qInv = qtInvert( q );\n" " float4 vcpy = vec;\n" " vcpy.w = 0.f;\n" -" float4 out = qtMul(qtMul(q,vcpy),qInv);\n" +" float4 out = qtMul(q,vcpy);\n" +" out = qtMul(out,qInv);\n" " return out;\n" "}\n" "\n" "__inline\n" -"Quaternion qtInvert(Quaternion q)\n" +" Quaternion qtInvert(Quaternion q)\n" "{\n" " return (Quaternion)(-q.xyz, q.w);\n" "}\n" "\n" "__inline\n" -"float4 qtInvRotate(const Quaternion q, float4 vec)\n" +" float4 qtInvRotate(const Quaternion q, float4 vec)\n" "{\n" " return qtRotate( qtInvert( q ), vec );\n" "}\n" "\n" -"__inline\n" -"float4 transform(const float4* p, const float4* translation, const Quaternion* orientation)\n" -"{\n" -" return qtRotate( *orientation, *p ) + (*translation);\n" -"}\n" +"\n" "\n" "void trInverse(float4 translationIn, Quaternion orientationIn,\n" -" float4* translationOut, Quaternion* orientationOut)\n" +" float4* translationOut, Quaternion* orientationOut)\n" "{\n" " *orientationOut = qtInvert(orientationIn);\n" " *translationOut = qtRotate(*orientationOut, -translationIn);\n" "}\n" "\n" -"void trMul(float4 translationA, Quaternion orientationA,\n" -" float4 translationB, Quaternion orientationB,\n" -" float4* translationOut, Quaternion* orientationOut)\n" -"{\n" -" *orientationOut = qtMul(orientationA,orientationB);\n" -" *translationOut = transform(&translationB,&translationA,&orientationA);\n" -"}\n" +"\n" "\n" "\n" "\n" @@ -173,9 +162,9 @@ static const char* rayCastKernelCL= \ "{\n" " rayFromLocal.w = 0.f;\n" " rayToLocal.w = 0.f;\n" -" bool result = true;\n" -" \n" -" float exitFraction = *hitFraction;\n" +" bool result = true;\n" +"\n" +" float exitFraction = hitFraction[0];\n" " float enterFraction = -0.3f;\n" " float4 curHitNormal = (float4)(0,0,0,0);\n" " for (int i=0;i 0.0)\n" -" {\n" -" float t = (-B - sqrt(D))/A;\n" +" if (D > 0.0f)\n" +" {\n" +" float t = (-B - sqrt(D))/A;\n" "\n" -" if ( (t >= 0.0f) && (t < (*hitFraction)) )\n" -" {\n" -" *hitFraction = t;\n" -" return true;\n" -" }\n" +" if ( (t >= 0.0f) && (t < (*hitFraction)) )\n" +" {\n" +" *hitFraction = t;\n" +" return true;\n" +" }\n" " }\n" " return false;\n" "}\n" "\n" "float4 setInterpolate3(float4 from, float4 to, float t)\n" "{\n" -" float s = 1.0f - t;\n" -" float4 result;\n" -" result = s * from + t * to;\n" -" result.w = 0.f; \n" -" return result; \n" +" float s = 1.0f - t;\n" +" float4 result;\n" +" result = s * from + t * to;\n" +" result.w = 0.f; \n" +" return result; \n" "}\n" "\n" "__kernel void rayCastKernel( \n" @@ -274,72 +266,76 @@ static const char* rayCastKernelCL= \ "{\n" "\n" " int i = get_global_id(0);\n" -" if (i=numRays)\n" +" return;\n" "\n" -" float4 rayFrom = rays[i].m_from;\n" -" float4 rayTo = rays[i].m_to;\n" -" float hitFraction = 1.f;\n" -" float4 hitPoint;\n" -" float4 hitNormal;\n" -" int hitBodyIndex= -1;\n" -" \n" -" \n" -" \n" -" \n" -" for (int b=0;b=0)\n" +" if (cachedCollidable.m_shapeType == SHAPE_CONVEX_HULL)\n" " {\n" -" hitResults[i].m_hitFraction = hitFraction;\n" -" hitResults[i].m_hitPoint = hitPoint;\n" -" hitResults[i].m_hitNormal = normalize(hitNormal);\n" -" hitResults[i].m_hitResult0 = hitBodyIndex;\n" +"\n" +" float4 invPos = (float4)(0,0,0,0);\n" +" float4 invOrn = (float4)(0,0,0,0);\n" +" float4 rayFromLocal = (float4)(0,0,0,0);\n" +" float4 rayToLocal = (float4)(0,0,0,0);\n" +" invOrn = qtInvert(orn);\n" +" invPos = qtRotate(invOrn, -pos);\n" +" rayFromLocal = qtRotate( invOrn, rayFrom ) + invPos;\n" +" rayToLocal = qtRotate( invOrn, rayTo) + invPos;\n" +" rayFromLocal.w = 0.f;\n" +" rayToLocal.w = 0.f;\n" +" int numFaces = convexShapes[cachedCollidable.m_shapeIndex].m_numFaces;\n" +" int faceOffset = convexShapes[cachedCollidable.m_shapeIndex].m_faceOffset;\n" +" if (numFaces)\n" +" {\n" +" if (rayConvex(rayFromLocal, rayToLocal, numFaces, faceOffset,faces, &hitFraction, &hitNormal))\n" +" {\n" +" hitBodyIndex = b;\n" +" }\n" +" }\n" +" }\n" +" if (cachedCollidable.m_shapeType == SHAPE_SPHERE)\n" +" {\n" +" float radius = cachedCollidable.m_radius;\n" +" \n" +" if (sphere_intersect(pos, radius, rayFrom, rayTo, &hitFraction))\n" +" {\n" +" hitBodyIndex = b;\n" +" hitPoint = setInterpolate3(rayFrom, rayTo,hitFraction);\n" +" hitNormal = (float4) (hitPoint-bodies[b].m_pos);\n" +" }\n" " }\n" " }\n" +"\n" +" if (hitBodyIndex>=0)\n" +" {\n" +" hitResults[i].m_hitFraction = hitFraction;\n" +" hitResults[i].m_hitPoint = hitPoint;\n" +" hitResults[i].m_hitNormal = normalize(hitNormal);\n" +" hitResults[i].m_hitResult0 = hitBodyIndex;\n" +" }\n" +"\n" "}\n" "\n" ;