From f10eb86f55e8d219956b247a26ae038a34f7875b Mon Sep 17 00:00:00 2001 From: erwin coumans Date: Wed, 19 Jun 2013 00:12:06 -0700 Subject: [PATCH] enable the GPU ray test in BasicGpuDemo, only for spheres at the moment --- .../Raycast/kernels/rayCastKernels.cl | 95 ++++++++++++++----- .../Raycast/kernels/rayCastKernels.h | 95 ++++++++++++++----- .../RigidBody/b3GpuRigidBodyPipeline.cpp | 2 +- 3 files changed, 139 insertions(+), 53 deletions(-) diff --git a/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.cl b/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.cl index a1a2dd744..9a627c141 100644 --- a/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.cl +++ b/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.cl @@ -1,4 +1,11 @@ +#define SHAPE_CONVEX_HULL 3 +#define SHAPE_PLANE 4 +#define SHAPE_CONCAVE_TRIMESH 5 +#define SHAPE_COMPOUND_OF_CONVEX_HULLS 6 +#define SHAPE_SPHERE 7 + + typedef struct { float4 m_from; @@ -30,67 +37,103 @@ typedef struct typedef struct Collidable { - int m_unused1; - int m_unused2; + union { + int m_numChildShapes; + int m_bvhIndex; + }; + float m_radius; int m_shapeType; int m_shapeIndex; } Collidable; -bool sphere_intersect(float4 spherePos, float radius, float4 rayFrom, float4 rayTo) + + +bool sphere_intersect(float4 spherePos, float radius, float4 rayFrom, float4 rayTo, float* hitFraction) { - // rs = ray.org - sphere.center - float4 rs = rayFrom - spherePos; - rs.w = 0.f; - float4 rayDir = (rayTo-rayFrom); + float4 rs = rayFrom - spherePos; + rs.w = 0.f; + float4 rayDir = rayTo-rayFrom; rayDir.w = 0.f; - rayDir = normalize(rayDir); + float A = dot(rayDir,rayDir); + float B = dot(rs, rayDir); + float C = dot(rs, rs) - (radius * radius); + + float D = B * B - A*C; - float B = dot(rs, rayDir); - float C = dot(rs, rs) - (radius * radius); - float D = B * B - C; + if (D > 0.0) + { + float t = (-B - sqrt(D))/A; - if (D > 0.0) - { - float t = -B - sqrt(D); - if ( (t > 0.0))// && (t < isect.t) ) + if ( (t >= 0.0f) && (t < (*hitFraction)) ) { - return true;//isect.t = t; - } + *hitFraction = t; + return true; + } } return false; } +float4 setInterpolate3(float4 from, float4 to, float t) +{ + float s = 1.0f - t; + float4 result; + result = s * from + t * to; + result.w = 0.f; + return result; +} + __kernel void rayCastKernel( int numRays, const __global b3RayInfo* rays, - __global b3RayHit* hits, + __global b3RayHit* hitResults, const int numBodies, __global Body* bodies, __global Collidable* collidables) { - bool hit=false; int i = get_global_id(0); if (i=0) + { + hitResults[i].m_hitFraction = hitFraction; + hitResults[i].m_hitPoint = setInterpolate3(rayFrom, rayTo,hitFraction); + float4 hitNormal = (float4) (hitResults[i].m_hitPoint-bodies[hitBodyIndex].m_pos); + hitResults[i].m_hitNormal = normalize(hitNormal); + hitResults[i].m_hitResult0 = hitBodyIndex; } - if (hit) - hits[i].m_hitFraction = 0.f; } } diff --git a/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h b/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h index 4a0e9e478..3d1268c19 100644 --- a/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h +++ b/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h @@ -1,6 +1,13 @@ //this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project static const char* rayCastKernelCL= \ "\n" +"#define SHAPE_CONVEX_HULL 3\n" +"#define SHAPE_PLANE 4\n" +"#define SHAPE_CONCAVE_TRIMESH 5\n" +"#define SHAPE_COMPOUND_OF_CONVEX_HULLS 6\n" +"#define SHAPE_SPHERE 7\n" +"\n" +"\n" "typedef struct\n" "{\n" " float4 m_from;\n" @@ -32,68 +39,104 @@ static const char* rayCastKernelCL= \ "\n" "typedef struct Collidable\n" "{\n" -" int m_unused1;\n" -" int m_unused2;\n" +" union {\n" +" int m_numChildShapes;\n" +" int m_bvhIndex;\n" +" };\n" +" float m_radius;\n" " int m_shapeType;\n" " int m_shapeIndex;\n" "} Collidable;\n" "\n" -"bool sphere_intersect(float4 spherePos, float radius, float4 rayFrom, float4 rayTo)\n" +"\n" +"\n" +"bool sphere_intersect(float4 spherePos, float radius, float4 rayFrom, float4 rayTo, float* hitFraction)\n" "{\n" -" // rs = ray.org - sphere.center\n" -" float4 rs = rayFrom - spherePos;\n" -" rs.w = 0.f;\n" -" float4 rayDir = (rayTo-rayFrom);\n" +" float4 rs = rayFrom - spherePos;\n" +" rs.w = 0.f;\n" +" float4 rayDir = rayTo-rayFrom;\n" " rayDir.w = 0.f;\n" -" rayDir = normalize(rayDir);\n" +" float A = dot(rayDir,rayDir);\n" +" float B = dot(rs, rayDir);\n" +" float C = dot(rs, rs) - (radius * radius);\n" +" \n" +" float D = B * B - A*C;\n" "\n" -" float B = dot(rs, rayDir);\n" -" float C = dot(rs, rs) - (radius * radius);\n" -" float D = B * B - C;\n" +" if (D > 0.0)\n" +" {\n" +" float t = (-B - sqrt(D))/A;\n" "\n" -" if (D > 0.0)\n" -" {\n" -" float t = -B - sqrt(D);\n" -" if ( (t > 0.0))// && (t < isect.t) )\n" +" if ( (t >= 0.0f) && (t < (*hitFraction)) )\n" " {\n" -" return true;//isect.t = t;\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" +"}\n" +"\n" "__kernel void rayCastKernel( \n" " int numRays, \n" " const __global b3RayInfo* rays, \n" -" __global b3RayHit* hits, \n" +" __global b3RayHit* hitResults, \n" " const int numBodies, \n" " __global Body* bodies,\n" " __global Collidable* collidables)\n" "{\n" "\n" -" bool hit=false;\n" "\n" " int i = get_global_id(0);\n" " if (i=0)\n" +" {\n" +" hitResults[i].m_hitFraction = hitFraction;\n" +" hitResults[i].m_hitPoint = setInterpolate3(rayFrom, rayTo,hitFraction);\n" +" float4 hitNormal = (float4) (hitResults[i].m_hitPoint-bodies[hitBodyIndex].m_pos);\n" +" hitResults[i].m_hitNormal = normalize(hitNormal);\n" +" hitResults[i].m_hitResult0 = hitBodyIndex;\n" " }\n" -" if (hit)\n" -" hits[i].m_hitFraction = 0.f;\n" " }\n" "}\n" "\n" diff --git a/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.cpp b/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.cpp index 29f4358ab..e00f058a3 100644 --- a/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.cpp +++ b/src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.cpp @@ -458,6 +458,6 @@ int b3GpuRigidBodyPipeline::registerPhysicsInstance(float mass, const float* po void b3GpuRigidBodyPipeline::castRays(const b3AlignedObjectArray& rays, b3AlignedObjectArray& hitResults) { - this->m_data->m_raycaster->castRaysHost(rays,hitResults,getNumBodies(),this->m_data->m_narrowphase->getBodiesCpu(),m_data->m_narrowphase->getNumCollidablesGpu(), m_data->m_narrowphase->getCollidablesCpu()); + this->m_data->m_raycaster->castRays(rays,hitResults,getNumBodies(),this->m_data->m_narrowphase->getBodiesCpu(),m_data->m_narrowphase->getNumCollidablesGpu(), m_data->m_narrowphase->getCollidablesCpu()); }