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

add missing rayCast kernel

Browse Source
This commit is contained in:
erwin coumans
2013-06-07 20:47:47 -07:00
parent afecad3ee4
commit 3f10082aa3
6 changed files with 119 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
Demos3/GpuDemos/main_opengl3core.cpp
View File

@@ -409,7 +409,7 @@ void myprintf(const char* msg)
int main(int argc, char* argv[]) int main(int argc, char* argv[])
{ {
b3OpenCLUtils::setCachePath("/Users/erwincoumans/develop/mycache"); //b3OpenCLUtils::setCachePath("/Users/erwincoumans/develop/mycache");
b3SetCustomPrintfFunc(myprintf); b3SetCustomPrintfFunc(myprintf);
b3Vector3 test(1,2,3); b3Vector3 test(1,2,3);

19
build/stringify.bat
View File

@@ -16,15 +16,16 @@ premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/Narrowpha
premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.cl" --headerfile="../src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.h" --stringname="bvhTraversalKernelCL" stringify premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.cl" --headerfile="../src/Bullet3OpenCL/NarrowphaseCollision/kernels/bvhTraversal.h" --stringname="bvhTraversalKernelCL" stringify
premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/integrateKernel.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/integrateKernel.h" --stringname="integrateKernelCL" stringify premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/integrateKernel.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/integrateKernel.h" --stringname="integrateKernelCL" stringify
premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/updateAabbsKernel.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/updateAabbsKernel.h" --stringname="updateAabbsKernelCL" stringify premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/updateAabbsKernel.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/updateAabbsKernel.h" --stringname="updateAabbsKernelCL" stringify
premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/solverSetup.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/solverSetup.h" --stringname="solverSetupCL" stringify premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/solverSetup.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/solverSetup.h" --stringname="solverSetupCL" stringify
premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/solverSetup2.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/solverSetup2.h" --stringname="solverSetup2CL" stringify premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/solverSetup2.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/solverSetup2.h" --stringname="solverSetup2CL" stringify
premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/batchingKernels.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/batchingKernels.h" --stringname="batchingKernelsCL" stringify premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/batchingKernels.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/batchingKernels.h" --stringname="batchingKernelsCL" stringify
premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/batchingKernelsNew.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/batchingKernelsNew.h" --stringname="batchingKernelsNewCL" stringify premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/batchingKernelsNew.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/batchingKernelsNew.h" --stringname="batchingKernelsNewCL" stringify
premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/solverUtils.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/solverUtils.h" --stringname="solverUtilsCL" stringify premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/solverUtils.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/solverUtils.h" --stringname="solverUtilsCL" stringify
premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/solveContact.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/solveContact.h" --stringname="solveContactCL" stringify premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/solveContact.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/solveContact.h" --stringname="solveContactCL" stringify
premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody//kernels/solveFriction.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/solveFriction.h" --stringname="solveFrictionCL" stringify premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/RigidBody/kernels/solveFriction.cl" --headerfile="../src/Bullet3OpenCL/RigidBody//kernels/solveFriction.h" --stringname="solveFrictionCL" stringify
premake4 --file=stringifyKernel.lua --kernelfile="../src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.cl" --headerfile="../src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h" --stringname="rayCastKernelCL" stringify
pause pause

5
src/Bullet3OpenCL/Initialize/b3OpenCLUtils.cpp
View File

@@ -735,8 +735,9 @@ cl_program b3OpenCLUtils_compileCLProgramFromString(cl_context clContext, cl_dev
} }
//we should make sure the src file exists so we can verify the timestamp with binary //we should make sure the src file exists so we can verify the timestamp with binary
assert(0); // assert(0);
fileUpToDate = false; b3Warning("Warning: cannot find OpenCL kernel %s to verify timestamp of binary cached kernel %s\n",clFileNameForCaching, binaryFileName);
fileUpToDate = true;
#else #else
//if we cannot find the source, assume it is OK in release builds //if we cannot find the source, assume it is OK in release builds
fileUpToDate = true; fileUpToDate = true;

3
src/Bullet3OpenCL/Raycast/b3GpuRaycast.cpp
View File

@@ -6,7 +6,7 @@
#include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h" #include "Bullet3OpenCL/Initialize/b3OpenCLUtils.h"
#include "Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h" #include "Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h"
#include "Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h" #include "Bullet3OpenCL/ParallelPrimitives/b3LauncherCL.h"
#include "Bullet3OpenCL/Raycast/kernels/rayCastKernels.h"
#define B3_RAYCAST_PATH "src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.cl" #define B3_RAYCAST_PATH "src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.cl"
@@ -30,7 +30,6 @@ b3GpuRaycast::b3GpuRaycast(cl_context ctx,cl_device_id device, cl_command_queue
m_data->m_q = q; m_data->m_q = q;
m_data->m_raytraceKernel = 0; m_data->m_raytraceKernel = 0;
char* rayCastKernelCL= 0;
{ {
cl_int errNum=0; cl_int errNum=0;

100
src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h Normal file
View File

@@ -0,0 +1,100 @@
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
static const char* rayCastKernelCL= \
"\n"
"typedef struct\n"
"{\n"
" float4 m_from;\n"
" float4 m_to;\n"
"} b3RayInfo;\n"
"\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"
"} b3RayHit;\n"
"\n"
"typedef struct\n"
"{\n"
" float4 m_pos;\n"
" float4 m_quat;\n"
" float4 m_linVel;\n"
" float4 m_angVel;\n"
"\n"
" unsigned int m_collidableIdx;\n"
" float m_invMass;\n"
" float m_restituitionCoeff;\n"
" float m_frictionCoeff;\n"
"} Body;\n"
"\n"
"typedef struct Collidable\n"
"{\n"
" int m_unused1;\n"
" int m_unused2;\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"
" // rs = ray.org - sphere.center\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"
"\n"
" float B = dot(rs, rayDir);\n"
" float C = dot(rs, rs) - (radius * radius);\n"
" float D = B * B - C;\n"
"\n"
" if (D > 0.0)\n"
" {\n"
" float t = -B - sqrt(D);\n"
" if ( (t > 0.0))// && (t < isect.t) )\n"
" {\n"
" return true;//isect.t = t;\n"
" }\n"
" }\n"
" return false;\n"
"}\n"
"\n"
"__kernel void rayCastKernel( \n"
" int numRays, \n"
" const __global b3RayInfo* rays, \n"
" __global b3RayHit* hits, \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"
"\n"
" float4 rayFrom = rays[i].m_from;\n"
" float4 rayTo = rays[i].m_to;\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"
" \n"
" float radius = 1.f;\n"
" \n"
" if (sphere_intersect(pos, radius, rayFrom, rayTo))\n"
" hit = true;\n"
" }\n"
" if (hit)\n"
" hits[i].m_hitFraction = 0.f;\n"
" }\n"
"}\n"
"\n"
;

4
src/Bullet3OpenCL/RigidBody/b3Config.h
View File

@@ -19,7 +19,11 @@ struct b3Config
int m_maxTriConvexPairCapacity; int m_maxTriConvexPairCapacity;
b3Config() b3Config()
#ifdef __APPLE__
:m_maxConvexBodies(64*1024), :m_maxConvexBodies(64*1024),
#else
:m_maxConvexBodies(128*1024),
#endif
m_maxConvexShapes(8192), m_maxConvexShapes(8192),
m_maxVerticesPerFace(64), m_maxVerticesPerFace(64),
m_maxFacesPerShape(64), m_maxFacesPerShape(64),