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

enable the GPU ray test in BasicGpuDemo, only for spheres at the moment

Browse Source
This commit is contained in:
erwin coumans
2013-06-19 00:12:06 -07:00
parent 481d54967f
commit f10eb86f55
3 changed files with 139 additions and 53 deletions
Download Patch File Download Diff File Expand all files Collapse all files

95
src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.cl
View File

@@ -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<numRays)
{
hits[i].m_hitFraction = 1.f;
hitResults[i].m_hitFraction = 1.f;
float4 rayFrom = rays[i].m_from;
float4 rayTo = rays[i].m_to;
float hitFraction = 1.f;
int hitBodyIndex= -1;
int cachedCollidableIndex = -1;
Collidable cachedCollidable;
for (int b=0;b<numBodies;b++)
{
float4 pos = bodies[b].m_pos;
// float4 orn = bodies[b].m_quat;
if (cachedCollidableIndex !=bodies[b].m_collidableIdx)
{
cachedCollidableIndex = bodies[b].m_collidableIdx;
cachedCollidable = collidables[cachedCollidableIndex];
}
float radius = 1.f;
if (sphere_intersect(pos, radius, rayFrom, rayTo))
hit = true;
if (cachedCollidable.m_shapeType == SHAPE_SPHERE)
{
float radius = cachedCollidable.m_radius;
if (sphere_intersect(pos, radius, rayFrom, rayTo, &hitFraction))
{
hitBodyIndex = b;
}
}
}
if (hitBodyIndex>=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;
}
}

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

@@ -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<numRays)\n"
" {\n"
" hits[i].m_hitFraction = 1.f;\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"
" int hitBodyIndex= -1;\n"
" \n"
" int cachedCollidableIndex = -1; \n"
" Collidable cachedCollidable;\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"
" cachedCollidableIndex = bodies[b].m_collidableIdx;\n"
" cachedCollidable = collidables[cachedCollidableIndex];\n"
" }\n"
" \n"
" float radius = 1.f;\n"
" \n"
" if (sphere_intersect(pos, radius, rayFrom, rayTo))\n"
" hit = true;\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"
" }\n"
" }\n"
" }\n"
" \n"
" if (hitBodyIndex>=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"

2
src/Bullet3OpenCL/RigidBody/b3GpuRigidBodyPipeline.cpp
View File

@@ -458,6 +458,6 @@ int b3GpuRigidBodyPipeline::registerPhysicsInstance(float mass, const float* po
void b3GpuRigidBodyPipeline::castRays(const b3AlignedObjectArray<b3RayInfo>& rays, b3AlignedObjectArray<b3RayHit>& 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());
}