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

fix ray-convex header file

Browse Source
This commit is contained in:
erwincoumans
2013-06-20 16:41:11 -07:00
parent d935eb2520
commit e39af46403
2 changed files with 119 additions and 123 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp
View File

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

238
src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h
View File

@@ -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<numFaces && result;i++)\n"
@@ -212,13 +201,16 @@ static const char* rayCastKernelCL= \
" if (exitFraction <= enterFraction)\n"
" result = false;\n"
" }\n"
" \n"
" result = result && (enterFraction < 0.f);\n"
" \n"
"\n"
" if (enterFraction < 0.f)\n"
" {\n"
" result = false;\n"
" }\n"
"\n"
" if (result)\n"
" { \n"
" *hitFraction = enterFraction;\n"
" *hitNormal = curHitNormal;\n"
" hitFraction[0] = enterFraction;\n"
" hitNormal[0] = curHitNormal;\n"
" }\n"
" return result;\n"
"}\n"
@@ -237,29 +229,29 @@ static const char* rayCastKernelCL= \
" float A = dot(rayDir,rayDir);\n"
" float B = dot(rs, rayDir);\n"
" float C = dot(rs, rs) - (radius * radius);\n"
" \n"
"\n"
" float D = B * B - A*C;\n"
"\n"
" if (D > 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"
" {\n"
" hitResults[i].m_hitFraction = 1.f;\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<numBodies;b++)\n"
" hitResults[i].m_hitFraction = 1.f;\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"
" int cachedCollidableIndex = -1;\n"
" Collidable cachedCollidable;\n"
"\n"
" for (int b=0;b<numBodies;b++)\n"
" {\n"
"\n"
" Body body = bodies[b];\n"
" float4 pos = body.m_pos;\n"
" float4 orn = body.m_quat;\n"
" if (cachedCollidableIndex != body.m_collidableIdx)\n"
" {\n"
" \n"
" float4 pos = bodies[b].m_pos;\n"
" float4 orn = bodies[b].m_quat;\n"
" \n"
" int cachedCollidableIndex = bodies[b].m_collidableIdx; \n"
" Collidable cachedCollidable = collidables[cachedCollidableIndex];\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"
" if (cachedCollidable.m_shapeType == SHAPE_CONVEX_HULL)\n"
" {\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"
" \n"
" trInverse(pos,orn, &invPos, &invOrn);\n"
" rayFromLocal = transform(&rayFrom, &invPos, &invOrn);\n"
" rayToLocal = transform(&rayTo, &invPos, &invOrn);\n"
" \n"
" int numFaces = convexShapes[cachedCollidable.m_shapeIndex].m_numFaces;\n"
" int faceOffset = convexShapes[cachedCollidable.m_shapeIndex].m_faceOffset;\n"
" \n"
" if (rayConvex(rayFromLocal, rayToLocal, numFaces, faceOffset,faces, &hitFraction, &hitNormal))\n"
" {\n"
" hitBodyIndex = b;\n"
" }\n"
" }\n"
" \n"
" cachedCollidableIndex = body.m_collidableIdx;\n"
" cachedCollidable = collidables[cachedCollidableIndex];\n"
" }\n"
" \n"
" if (hitBodyIndex>=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"
;