diff --git a/Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp b/Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp index 83f9c4d79..b0989a2e8 100644 --- a/Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp +++ b/Demos3/GpuDemos/rigidbody/GpuConvexScene.cpp @@ -234,8 +234,8 @@ GpuRaytraceScene::GpuRaytraceScene() m_raytraceData = new GpuRaytraceInternalData; m_raytraceData->m_texId = new GLuint; - m_raytraceData->textureWidth = 1024;//1024; - m_raytraceData->textureHeight = 1024; + m_raytraceData->textureWidth = 256;//1024;//1024; + m_raytraceData->textureHeight = 256;//1024; //create new texture glGenTextures(1, m_raytraceData->m_texId); @@ -471,9 +471,9 @@ void GpuRaytraceScene::renderScene() if (hit) { - m_raytraceData->m_texels[(i)*3+0] = 255; - m_raytraceData->m_texels[(i)*3+1] = 0; - m_raytraceData->m_texels[(i)*3+2] = 0; + m_raytraceData->m_texels[(i)*3+0] = 128+128.f*hits[i].m_hitNormal.x; + m_raytraceData->m_texels[(i)*3+1] = 128+128.f*hits[i].m_hitNormal.y; + m_raytraceData->m_texels[(i)*3+2] = 128+128.f*hits[i].m_hitNormal.z; } else { m_raytraceData->m_texels[(i)*3+0] = 0; diff --git a/src/Bullet3OpenCL/Initialize/b3OpenCLUtils.cpp b/src/Bullet3OpenCL/Initialize/b3OpenCLUtils.cpp index ed45621d9..1d305544f 100644 --- a/src/Bullet3OpenCL/Initialize/b3OpenCLUtils.cpp +++ b/src/Bullet3OpenCL/Initialize/b3OpenCLUtils.cpp @@ -610,7 +610,7 @@ cl_program b3OpenCLUtils_compileCLProgramFromString(cl_context clContext, cl_dev int fileUpToDate = 0; int binaryFileValid=0; - if (clFileNameForCaching) + if (!disableBinaryCaching && clFileNameForCaching) { clGetDeviceInfo(device, CL_DEVICE_NAME, 256, &deviceName, NULL); clGetDeviceInfo(device, CL_DRIVER_VERSION, 256, &driverVersion, NULL); @@ -918,7 +918,7 @@ cl_program b3OpenCLUtils_compileCLProgramFromString(cl_context clContext, cl_dev } - if( clFileNameForCaching ) + if( !disableBinaryCaching && clFileNameForCaching ) { // write to binary cl_uint numAssociatedDevices; diff --git a/src/Bullet3OpenCL/Raycast/b3GpuRaycast.cpp b/src/Bullet3OpenCL/Raycast/b3GpuRaycast.cpp index 9dff7fc89..f71aa7795 100644 --- a/src/Bullet3OpenCL/Raycast/b3GpuRaycast.cpp +++ b/src/Bullet3OpenCL/Raycast/b3GpuRaycast.cpp @@ -209,6 +209,8 @@ void b3GpuRaycast::castRays(const b3AlignedObjectArray& rays, b3Align int numBodies,const struct b3RigidBodyCL* bodies, int numCollidables, const struct b3Collidable* collidables, const struct b3GpuNarrowPhaseInternalData* narrowphaseData) { + //castRaysHost(rays,hitResults,numBodies,bodies,numCollidables,collidables,narrowphaseData); + B3_PROFILE("castRaysGPU"); b3OpenCLArray gpuRays(m_data->m_context,m_data->m_q); diff --git a/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.cl b/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.cl index 8bc21f51b..95a3f21dc 100644 --- a/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.cl +++ b/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.cl @@ -14,12 +14,12 @@ typedef struct typedef struct { - float m_hitFraction; - int m_hitResult0; - int m_hitResult1; - int m_hitResult2; - float4 m_hitPoint; - float4 m_hitNormal; + float m_hitFraction; + int m_hitResult0; + int m_hitResult1; + int m_hitResult2; + float4 m_hitPoint; + float4 m_hitNormal; } b3RayHit; typedef struct @@ -53,12 +53,12 @@ typedef struct float4 m_extents; float4 mC; float4 mE; - + float m_radius; int m_faceOffset; int m_numFaces; int m_numVertices; - + int m_vertexOffset; int m_uniqueEdgesOffset; int m_numUniqueEdges; @@ -82,20 +82,18 @@ typedef struct typedef float4 Quaternion; __inline -Quaternion qtMul(Quaternion a, Quaternion b); + Quaternion qtMul(Quaternion a, Quaternion b); __inline -Quaternion qtNormalize(Quaternion in); - -__inline -float4 qtRotate(Quaternion q, float4 vec); - -__inline -Quaternion qtInvert(Quaternion q); + Quaternion qtNormalize(Quaternion in); __inline -float dot3F4(float4 a, float4 b) + Quaternion qtInvert(Quaternion q); + + +__inline + float dot3F4(float4 a, float4 b) { float4 a1 = (float4)(a.xyz,0.f); float4 b1 = (float4)(b.xyz,0.f); @@ -104,65 +102,56 @@ float dot3F4(float4 a, float4 b) __inline -Quaternion qtMul(Quaternion a, Quaternion b) + Quaternion qtMul(Quaternion a, Quaternion b) { Quaternion ans; ans = cross( a, b ); ans += a.w*b+b.w*a; -// ans.w = a.w*b.w - (a.x*b.x+a.y*b.y+a.z*b.z); + // ans.w = a.w*b.w - (a.x*b.x+a.y*b.y+a.z*b.z); ans.w = a.w*b.w - dot3F4(a, b); return ans; } __inline -Quaternion qtNormalize(Quaternion in) + Quaternion qtNormalize(Quaternion in) { return fast_normalize(in); -// in /= length( in ); -// return in; + // in /= length( in ); + // return in; } __inline -float4 qtRotate(Quaternion q, float4 vec) + float4 qtRotate(Quaternion q, float4 vec) { Quaternion qInv = qtInvert( q ); float4 vcpy = vec; vcpy.w = 0.f; - float4 out = qtMul(qtMul(q,vcpy),qInv); + float4 out = qtMul(q,vcpy); + out = qtMul(out,qInv); return out; } __inline -Quaternion qtInvert(Quaternion q) + Quaternion qtInvert(Quaternion q) { return (Quaternion)(-q.xyz, q.w); } __inline -float4 qtInvRotate(const Quaternion q, float4 vec) + float4 qtInvRotate(const Quaternion q, float4 vec) { return qtRotate( qtInvert( q ), vec ); } -__inline -float4 transform(const float4* p, const float4* translation, const Quaternion* orientation) -{ - return qtRotate( *orientation, *p ) + (*translation); -} + void trInverse(float4 translationIn, Quaternion orientationIn, - float4* translationOut, Quaternion* orientationOut) + float4* translationOut, Quaternion* orientationOut) { *orientationOut = qtInvert(orientationIn); *translationOut = qtRotate(*orientationOut, -translationIn); } -void trMul(float4 translationA, Quaternion orientationA, - float4 translationB, Quaternion orientationB, - float4* translationOut, Quaternion* orientationOut) -{ - *orientationOut = qtMul(orientationA,orientationB); - *translationOut = transform(&translationB,&translationA,&orientationA); -} + @@ -171,10 +160,10 @@ bool rayConvex(float4 rayFromLocal, float4 rayToLocal, int numFaces, int faceOff { rayFromLocal.w = 0.f; rayToLocal.w = 0.f; - bool result = true; - - float exitFraction = *hitFraction; - float enterFraction = -0.1f; + bool result = true; + + float exitFraction = hitFraction[0]; + float enterFraction = -0.3f; float4 curHitNormal = (float4)(0,0,0,0); for (int i=0;i 0.0) - { - float t = (-B - sqrt(D))/A; + if (D > 0.0f) + { + float t = (-B - sqrt(D))/A; - if ( (t >= 0.0f) && (t < (*hitFraction)) ) - { - *hitFraction = t; - return true; - } + if ( (t >= 0.0f) && (t < (*hitFraction)) ) + { + *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; + float s = 1.0f - t; + float4 result; + result = s * from + t * to; + result.w = 0.f; + return result; } __kernel void rayCastKernel( @@ -272,72 +264,74 @@ __kernel void rayCastKernel( { int i = get_global_id(0); - if (i=numRays) + return; - float4 rayFrom = rays[i].m_from; - float4 rayTo = rays[i].m_to; - float hitFraction = 1.f; - float4 hitPoint; - float4 hitNormal; - int hitBodyIndex= -1; - - int cachedCollidableIndex = -1; - Collidable cachedCollidable; - - for (int b=0;b=0) + if (cachedCollidable.m_shapeType == SHAPE_CONVEX_HULL) { - hitResults[i].m_hitFraction = hitFraction; - hitResults[i].m_hitPoint = hitPoint; - hitResults[i].m_hitNormal = normalize(hitNormal); - hitResults[i].m_hitResult0 = hitBodyIndex; + + float4 invPos = (float4)(0,0,0,0); + float4 invOrn = (float4)(0,0,0,0); + float4 rayFromLocal = (float4)(0,0,0,0); + float4 rayToLocal = (float4)(0,0,0,0); + invOrn = qtInvert(orn); + invPos = qtRotate(invOrn, -pos); + rayFromLocal = qtRotate( invOrn, rayFrom ) + invPos; + rayToLocal = qtRotate( invOrn, rayTo) + invPos; + rayFromLocal.w = 0.f; + rayToLocal.w = 0.f; + int numFaces = convexShapes[cachedCollidable.m_shapeIndex].m_numFaces; + int faceOffset = convexShapes[cachedCollidable.m_shapeIndex].m_faceOffset; + if (numFaces) + { + if (rayConvex(rayFromLocal, rayToLocal, numFaces, faceOffset,faces, &hitFraction, &hitNormal)) + { + hitBodyIndex = b; + } + } + } + if (cachedCollidable.m_shapeType == SHAPE_SPHERE) + { + float radius = cachedCollidable.m_radius; + + if (sphere_intersect(pos, radius, rayFrom, rayTo, &hitFraction)) + { + hitBodyIndex = b; + hitPoint = setInterpolate3(rayFrom, rayTo,hitFraction); + hitNormal = (float4) (hitPoint-bodies[b].m_pos); + } } } + + if (hitBodyIndex>=0) + { + hitResults[i].m_hitFraction = hitFraction; + hitResults[i].m_hitPoint = hitPoint; + hitResults[i].m_hitNormal = normalize(hitNormal); + hitResults[i].m_hitResult0 = hitBodyIndex; + } + } diff --git a/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h b/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h index 5c2f65c89..02d8a43a4 100644 --- a/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h +++ b/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h @@ -176,7 +176,7 @@ static const char* rayCastKernelCL= \ " bool result = true;\n" " \n" " float exitFraction = *hitFraction;\n" -" float enterFraction = -0.1f;\n" +" float enterFraction = -0.3f;\n" " float4 curHitNormal = (float4)(0,0,0,0);\n" " for (int i=0;i