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

Merge branch 'master' of https://github.com/erwincoumans/bullet3

Browse Source
This commit is contained in:
erwincoumans
2013-06-19 12:52:12 -07:00
parent 41b49b6e18 aa1c2db35a
commit 13a4084adf
21 changed files with 1908 additions and 149 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -41,7 +41,7 @@ bool gDebugSkipLoadingBinary = false;
#include <assert.h>
#define b3Assert assert
#ifdef __APPLE__
#ifndef _WIN32
#include <sys/stat.h>
#endif
@@ -752,7 +752,6 @@ cl_program b3OpenCLUtils_compileCLProgramFromString(cl_context clContext, cl_dev
#else
fileUpToDate = true;
#ifdef __APPLE__
if (mkdir(sCachedBinaryPath,0777) == -1)
{
}
@@ -760,7 +759,6 @@ cl_program b3OpenCLUtils_compileCLProgramFromString(cl_context clContext, cl_dev
{
b3Printf("Succesfully created cache directory: %s\n", sCachedBinaryPath);
}
#endif
#endif //_WIN32
}
@@ -873,10 +871,9 @@ cl_program b3OpenCLUtils_compileCLProgramFromString(cl_context clContext, cl_dev
size_t program_length = kernelSource ? strlen(kernelSource) : 0;
#ifdef MAC //or __APPLE__?
char* flags = "-cl-mad-enable -DMAC -DGUID_ARG";
char* flags = "-cl-mad-enable -DMAC ";
#else
//const char* flags = "-DGUID_ARG= -fno-alias";
const char* flags = "-DGUID_ARG= ";
const char* flags = "";
#endif

34
src/Bullet3OpenCL/ParallelPrimitives/b3OpenCLArray.h
View File

@@ -4,7 +4,7 @@
#include "Bullet3Common/b3AlignedObjectArray.h"
#include "Bullet3OpenCL/Initialize/b3OpenCLInclude.h"
template <typename T>
template <typename T>
class b3OpenCLArray
{
size_t m_size;
@@ -59,9 +59,9 @@ public:
m_size = sizeInElements;
m_capacity = sizeInElements;
}
// we could enable this assignment, but need to make sure to avoid accidental deep copies
// b3OpenCLArray<T>& operator=(const b3AlignedObjectArray<T>& src)
// b3OpenCLArray<T>& operator=(const b3AlignedObjectArray<T>& src)
// {
// copyFromArray(src);
// return *this;
@@ -73,16 +73,16 @@ public:
return m_clBuffer;
}
virtual ~b3OpenCLArray()
{
deallocate();
m_size=0;
m_capacity=0;
}
B3_FORCE_INLINE bool push_back(const T& _Val,bool waitForCompletion=true)
{
{
bool result = true;
size_t sz = size();
if( sz == capacity() )
@@ -147,7 +147,7 @@ public:
}
B3_FORCE_INLINE size_t capacity() const
{
{
return m_capacity;
}
@@ -164,9 +164,9 @@ public:
//create a new OpenCL buffer
size_t memSizeInBytes = sizeof(T)*_Count;
cl_mem buf = clCreateBuffer(m_clContext, CL_MEM_READ_WRITE, memSizeInBytes, NULL, &ciErrNum);
b3Assert(ciErrNum==CL_SUCCESS);
if (ciErrNum!=CL_SUCCESS)
{
b3Error("OpenCL out-of-memory\n");
_Count = 0;
result = false;
}
@@ -191,7 +191,7 @@ public:
deallocate();
m_clBuffer = buf;
m_capacity = _Count;
} else
{
@@ -212,14 +212,14 @@ public:
b3Assert(m_clBuffer);
b3Assert(destination);
//likely some error, destination is same as source
b3Assert(m_clBuffer != destination);
b3Assert((firstElem+numElements)<=m_size);
cl_int status = 0;
b3Assert(numElements>0);
b3Assert(numElements<=m_size);
@@ -227,7 +227,7 @@ public:
size_t srcOffsetBytes = sizeof(T)*firstElem;
size_t dstOffsetInBytes = sizeof(T)*dstOffsetInElems;
status = clEnqueueCopyBuffer( m_commandQueue, m_clBuffer, destination,
status = clEnqueueCopyBuffer( m_commandQueue, m_clBuffer, destination,
srcOffsetBytes, dstOffsetInBytes, sizeof(T)*numElements, 0, 0, 0 );
b3Assert( status == CL_SUCCESS );
@@ -236,7 +236,7 @@ public:
void copyFromHost(const b3AlignedObjectArray<T>& srcArray, bool waitForCompletion=true)
{
size_t newSize = srcArray.size();
bool copyOldContents = false;
resize (newSize,copyOldContents);
if (newSize)
@@ -262,7 +262,7 @@ public:
b3Error("copyFromHostPointer invalid range\n");
}
}
void copyToHost(b3AlignedObjectArray<T>& destArray, bool waitForCompletion=true) const
{
@@ -274,7 +274,7 @@ public:
void copyToHostPointer(T* destPtr, size_t numElem, size_t srcFirstElem=0, bool waitForCompletion=true) const
{
b3Assert(numElem+srcFirstElem <= capacity());
if(numElem+srcFirstElem <= capacity())
{
cl_int status = 0;
@@ -289,7 +289,7 @@ public:
b3Error("copyToHostPointer invalid range\n");
}
}
void copyFromOpenCLArray(const b3OpenCLArray& src)
{
size_t newSize = src.size();

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

@@ -87,7 +87,7 @@ void b3GpuRaycast::castRaysHost(const b3AlignedObjectArray<b3RayInfo>& rays, b3A
b3Vector3 rayTo = rays[r].m_to;
float hitFraction = hitResults[r].m_hitFraction;
int sphereHit = -1;
int hitBodyIndex= -1;
for (int b=0;b<numBodies;b++)
{
@@ -95,20 +95,35 @@ void b3GpuRaycast::castRaysHost(const b3AlignedObjectArray<b3RayInfo>& rays, b3A
const b3Vector3& pos = bodies[b].m_pos;
const b3Quaternion& orn = bodies[b].m_quat;
b3Scalar radius = 1;
if (sphere_intersect(pos, radius, rayFrom, rayTo,hitFraction))
switch (collidables[bodies[b].m_collidableIdx].m_shapeType)
{
sphereHit = b;
case SHAPE_SPHERE:
{
b3Scalar radius = collidables[bodies[b].m_collidableIdx].m_radius;
if (sphere_intersect(pos, radius, rayFrom, rayTo,hitFraction))
{
hitBodyIndex = b;
}
}
default:
{
static bool once=true;
if (once)
{
once=false;
b3Warning("Raytest: unsupported shape type\n");
}
}
}
}
if (sphereHit>=0)
if (hitBodyIndex>=0)
{
hitResults[r].m_hitFraction = hitFraction;
hitResults[r].m_hitPoint.setInterpolate3(rays[r].m_from, rays[r].m_to,hitFraction);
hitResults[r].m_hitNormal = (hitResults[r].m_hitPoint-bodies[sphereHit].m_pos).normalize();
hitResults[r].m_hitResult0 = sphereHit;
hitResults[r].m_hitNormal = (hitResults[r].m_hitPoint-bodies[hitBodyIndex].m_pos).normalize();
hitResults[r].m_hitResult0 = hitBodyIndex;
}
}

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());
}

4
src/Bullet3OpenCL/RigidBody/b3Solver.cpp
View File

@@ -791,7 +791,7 @@ void b3Solver::convertToConstraints( const b3OpenCLArray<b3RigidBodyCL>* bodyBuf
int nContacts, const ConstraintCfg& cfg )
{
b3OpenCLArray<b3GpuConstraint4>* constraintNative =0;
contactCOut->resize(nContacts);
struct CB
{
int m_nContacts;
@@ -825,7 +825,7 @@ void b3Solver::convertToConstraints( const b3OpenCLArray<b3RigidBodyCL>* bodyBuf
}
contactCOut->resize(nContacts);
}
/*