2017-08-01 12:30:58 +00:00
//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
2019-01-03 13:26:51 +00:00
static const char * satKernelsCL =
" //keep this enum in sync with the CPU version (in btCollidable.h) \n "
" //written by Erwin Coumans \n "
" #define SHAPE_CONVEX_HULL 3 \n "
" #define SHAPE_CONCAVE_TRIMESH 5 \n "
" #define TRIANGLE_NUM_CONVEX_FACES 5 \n "
" #define SHAPE_COMPOUND_OF_CONVEX_HULLS 6 \n "
" #define B3_MAX_STACK_DEPTH 256 \n "
" typedef unsigned int u32; \n "
" ///keep this in sync with btCollidable.h \n "
" typedef struct \n "
" { \n "
" union { \n "
" int m_numChildShapes; \n "
" int m_bvhIndex; \n "
" }; \n "
" union \n "
" { \n "
" float m_radius; \n "
" int m_compoundBvhIndex; \n "
" }; \n "
" \n "
" int m_shapeType; \n "
" int m_shapeIndex; \n "
" \n "
" } btCollidableGpu; \n "
" #define MAX_NUM_PARTS_IN_BITS 10 \n "
" ///b3QuantizedBvhNode is a compressed aabb node, 16 bytes. \n "
" ///Node can be used for leafnode or internal node. Leafnodes can point to 32-bit triangle index (non-negative range). \n "
" typedef struct \n "
" { \n "
" //12 bytes \n "
" unsigned short int m_quantizedAabbMin[3]; \n "
" unsigned short int m_quantizedAabbMax[3]; \n "
" //4 bytes \n "
" int m_escapeIndexOrTriangleIndex; \n "
" } b3QuantizedBvhNode; \n "
" typedef struct \n "
" { \n "
" float4 m_aabbMin; \n "
" float4 m_aabbMax; \n "
" float4 m_quantization; \n "
" int m_numNodes; \n "
" int m_numSubTrees; \n "
" int m_nodeOffset; \n "
" int m_subTreeOffset; \n "
" } b3BvhInfo; \n "
" int getTriangleIndex(const b3QuantizedBvhNode* rootNode) \n "
" { \n "
" unsigned int x=0; \n "
" unsigned int y = (~(x&0))<<(31-MAX_NUM_PARTS_IN_BITS); \n "
" // Get only the lower bits where the triangle index is stored \n "
" return (rootNode->m_escapeIndexOrTriangleIndex&~(y)); \n "
" } \n "
" int getTriangleIndexGlobal(__global const b3QuantizedBvhNode* rootNode) \n "
" { \n "
" unsigned int x=0; \n "
" unsigned int y = (~(x&0))<<(31-MAX_NUM_PARTS_IN_BITS); \n "
" // Get only the lower bits where the triangle index is stored \n "
" return (rootNode->m_escapeIndexOrTriangleIndex&~(y)); \n "
" } \n "
" int isLeafNode(const b3QuantizedBvhNode* rootNode) \n "
" { \n "
" //skipindex is negative (internal node), triangleindex >=0 (leafnode) \n "
" return (rootNode->m_escapeIndexOrTriangleIndex >= 0)? 1 : 0; \n "
" } \n "
" int isLeafNodeGlobal(__global const b3QuantizedBvhNode* rootNode) \n "
" { \n "
" //skipindex is negative (internal node), triangleindex >=0 (leafnode) \n "
" return (rootNode->m_escapeIndexOrTriangleIndex >= 0)? 1 : 0; \n "
" } \n "
" \n "
" int getEscapeIndex(const b3QuantizedBvhNode* rootNode) \n "
" { \n "
" return -rootNode->m_escapeIndexOrTriangleIndex; \n "
" } \n "
" int getEscapeIndexGlobal(__global const b3QuantizedBvhNode* rootNode) \n "
" { \n "
" return -rootNode->m_escapeIndexOrTriangleIndex; \n "
" } \n "
" typedef struct \n "
" { \n "
" //12 bytes \n "
" unsigned short int m_quantizedAabbMin[3]; \n "
" unsigned short int m_quantizedAabbMax[3]; \n "
" //4 bytes, points to the root of the subtree \n "
" int m_rootNodeIndex; \n "
" //4 bytes \n "
" int m_subtreeSize; \n "
" int m_padding[3]; \n "
" } b3BvhSubtreeInfo; \n "
" typedef struct \n "
" { \n "
" float4 m_childPosition; \n "
" float4 m_childOrientation; \n "
" int m_shapeIndex; \n "
" int m_unused0; \n "
" int m_unused1; \n "
" int m_unused2; \n "
" } btGpuChildShape; \n "
" typedef struct \n "
" { \n "
" float4 m_pos; \n "
" float4 m_quat; \n "
" float4 m_linVel; \n "
" float4 m_angVel; \n "
" u32 m_collidableIdx; \n "
" float m_invMass; \n "
" float m_restituitionCoeff; \n "
" float m_frictionCoeff; \n "
" } BodyData; \n "
" typedef struct \n "
" { \n "
" float4 m_localCenter; \n "
" float4 m_extents; \n "
" float4 mC; \n "
" float4 mE; \n "
" \n "
" float m_radius; \n "
" int m_faceOffset; \n "
" int m_numFaces; \n "
" int m_numVertices; \n "
" int m_vertexOffset; \n "
" int m_uniqueEdgesOffset; \n "
" int m_numUniqueEdges; \n "
" int m_unused; \n "
" } ConvexPolyhedronCL; \n "
" typedef struct \n "
" { \n "
" union \n "
" { \n "
" float4 m_min; \n "
" float m_minElems[4]; \n "
" int m_minIndices[4]; \n "
" }; \n "
" union \n "
" { \n "
" float4 m_max; \n "
" float m_maxElems[4]; \n "
" int m_maxIndices[4]; \n "
" }; \n "
" } btAabbCL; \n "
" #ifndef B3_AABB_H \n "
" #define B3_AABB_H \n "
" #ifndef B3_FLOAT4_H \n "
" #define B3_FLOAT4_H \n "
" #ifndef B3_PLATFORM_DEFINITIONS_H \n "
" #define B3_PLATFORM_DEFINITIONS_H \n "
" struct MyTest \n "
" { \n "
" int bla; \n "
" }; \n "
" #ifdef __cplusplus \n "
" #else \n "
" //keep B3_LARGE_FLOAT*B3_LARGE_FLOAT < FLT_MAX \n "
" #define B3_LARGE_FLOAT 1e18f \n "
" #define B3_INFINITY 1e18f \n "
" #define b3Assert(a) \n "
" #define b3ConstArray(a) __global const a* \n "
" #define b3AtomicInc atomic_inc \n "
" #define b3AtomicAdd atomic_add \n "
" #define b3Fabs fabs \n "
" #define b3Sqrt native_sqrt \n "
" #define b3Sin native_sin \n "
" #define b3Cos native_cos \n "
" #define B3_STATIC \n "
" #endif \n "
" #endif \n "
" #ifdef __cplusplus \n "
" #else \n "
" typedef float4 b3Float4; \n "
" #define b3Float4ConstArg const b3Float4 \n "
" #define b3MakeFloat4 (float4) \n "
" float b3Dot3F4(b3Float4ConstArg v0,b3Float4ConstArg v1) \n "
" { \n "
" float4 a1 = b3MakeFloat4(v0.xyz,0.f); \n "
" float4 b1 = b3MakeFloat4(v1.xyz,0.f); \n "
" return dot(a1, b1); \n "
" } \n "
" b3Float4 b3Cross3(b3Float4ConstArg v0,b3Float4ConstArg v1) \n "
" { \n "
" float4 a1 = b3MakeFloat4(v0.xyz,0.f); \n "
" float4 b1 = b3MakeFloat4(v1.xyz,0.f); \n "
" return cross(a1, b1); \n "
" } \n "
" #define b3MinFloat4 min \n "
" #define b3MaxFloat4 max \n "
" #define b3Normalized(a) normalize(a) \n "
" #endif \n "
" \n "
" inline bool b3IsAlmostZero(b3Float4ConstArg v) \n "
" { \n "
" if(b3Fabs(v.x)>1e-6 || b3Fabs(v.y)>1e-6 || b3Fabs(v.z)>1e-6) \n "
" return false; \n "
" return true; \n "
" } \n "
" inline int b3MaxDot( b3Float4ConstArg vec, __global const b3Float4* vecArray, int vecLen, float* dotOut ) \n "
" { \n "
" float maxDot = -B3_INFINITY; \n "
" int i = 0; \n "
" int ptIndex = -1; \n "
" for( i = 0; i < vecLen; i++ ) \n "
" { \n "
" float dot = b3Dot3F4(vecArray[i],vec); \n "
" \n "
" if( dot > maxDot ) \n "
" { \n "
" maxDot = dot; \n "
" ptIndex = i; \n "
" } \n "
" } \n "
" b3Assert(ptIndex>=0); \n "
" if (ptIndex<0) \n "
" { \n "
" ptIndex = 0; \n "
" } \n "
" *dotOut = maxDot; \n "
" return ptIndex; \n "
" } \n "
" #endif //B3_FLOAT4_H \n "
" #ifndef B3_MAT3x3_H \n "
" #define B3_MAT3x3_H \n "
" #ifndef B3_QUAT_H \n "
" #define B3_QUAT_H \n "
" #ifndef B3_PLATFORM_DEFINITIONS_H \n "
" #ifdef __cplusplus \n "
" #else \n "
" #endif \n "
" #endif \n "
" #ifndef B3_FLOAT4_H \n "
" #ifdef __cplusplus \n "
" #else \n "
" #endif \n "
" #endif //B3_FLOAT4_H \n "
" #ifdef __cplusplus \n "
" #else \n "
" typedef float4 b3Quat; \n "
" #define b3QuatConstArg const b3Quat \n "
" \n "
" \n "
" inline float4 b3FastNormalize4(float4 v) \n "
" { \n "
" v = (float4)(v.xyz,0.f); \n "
" return fast_normalize(v); \n "
" } \n "
" \n "
" inline b3Quat b3QuatMul(b3Quat a, b3Quat b); \n "
" inline b3Quat b3QuatNormalized(b3QuatConstArg in); \n "
" inline b3Quat b3QuatRotate(b3QuatConstArg q, b3QuatConstArg vec); \n "
" inline b3Quat b3QuatInvert(b3QuatConstArg q); \n "
" inline b3Quat b3QuatInverse(b3QuatConstArg q); \n "
" inline b3Quat b3QuatMul(b3QuatConstArg a, b3QuatConstArg b) \n "
" { \n "
" b3Quat ans; \n "
" ans = b3Cross3( a, b ); \n "
" ans += a.w*b+b.w*a; \n "
" // ans.w = a.w*b.w - (a.x*b.x+a.y*b.y+a.z*b.z); \n "
" ans.w = a.w*b.w - b3Dot3F4(a, b); \n "
" return ans; \n "
" } \n "
" inline b3Quat b3QuatNormalized(b3QuatConstArg in) \n "
" { \n "
" b3Quat q; \n "
" q=in; \n "
" //return b3FastNormalize4(in); \n "
" float len = native_sqrt(dot(q, q)); \n "
" if(len > 0.f) \n "
" { \n "
" q *= 1.f / len; \n "
" } \n "
" else \n "
" { \n "
" q.x = q.y = q.z = 0.f; \n "
" q.w = 1.f; \n "
" } \n "
" return q; \n "
" } \n "
" inline float4 b3QuatRotate(b3QuatConstArg q, b3QuatConstArg vec) \n "
" { \n "
" b3Quat qInv = b3QuatInvert( q ); \n "
" float4 vcpy = vec; \n "
" vcpy.w = 0.f; \n "
" float4 out = b3QuatMul(b3QuatMul(q,vcpy),qInv); \n "
" return out; \n "
" } \n "
" inline b3Quat b3QuatInverse(b3QuatConstArg q) \n "
" { \n "
" return (b3Quat)(-q.xyz, q.w); \n "
" } \n "
" inline b3Quat b3QuatInvert(b3QuatConstArg q) \n "
" { \n "
" return (b3Quat)(-q.xyz, q.w); \n "
" } \n "
" inline float4 b3QuatInvRotate(b3QuatConstArg q, b3QuatConstArg vec) \n "
" { \n "
" return b3QuatRotate( b3QuatInvert( q ), vec ); \n "
" } \n "
" inline b3Float4 b3TransformPoint(b3Float4ConstArg point, b3Float4ConstArg translation, b3QuatConstArg orientation) \n "
" { \n "
" return b3QuatRotate( orientation, point ) + (translation); \n "
" } \n "
" \n "
" #endif \n "
" #endif //B3_QUAT_H \n "
" #ifdef __cplusplus \n "
" #else \n "
" typedef struct \n "
" { \n "
" b3Float4 m_row[3]; \n "
" }b3Mat3x3; \n "
" #define b3Mat3x3ConstArg const b3Mat3x3 \n "
" #define b3GetRow(m,row) (m.m_row[row]) \n "
" inline b3Mat3x3 b3QuatGetRotationMatrix(b3Quat quat) \n "
" { \n "
" b3Float4 quat2 = (b3Float4)(quat.x*quat.x, quat.y*quat.y, quat.z*quat.z, 0.f); \n "
" b3Mat3x3 out; \n "
" out.m_row[0].x=1-2*quat2.y-2*quat2.z; \n "
" out.m_row[0].y=2*quat.x*quat.y-2*quat.w*quat.z; \n "
" out.m_row[0].z=2*quat.x*quat.z+2*quat.w*quat.y; \n "
" out.m_row[0].w = 0.f; \n "
" out.m_row[1].x=2*quat.x*quat.y+2*quat.w*quat.z; \n "
" out.m_row[1].y=1-2*quat2.x-2*quat2.z; \n "
" out.m_row[1].z=2*quat.y*quat.z-2*quat.w*quat.x; \n "
" out.m_row[1].w = 0.f; \n "
" out.m_row[2].x=2*quat.x*quat.z-2*quat.w*quat.y; \n "
" out.m_row[2].y=2*quat.y*quat.z+2*quat.w*quat.x; \n "
" out.m_row[2].z=1-2*quat2.x-2*quat2.y; \n "
" out.m_row[2].w = 0.f; \n "
" return out; \n "
" } \n "
" inline b3Mat3x3 b3AbsoluteMat3x3(b3Mat3x3ConstArg matIn) \n "
" { \n "
" b3Mat3x3 out; \n "
" out.m_row[0] = fabs(matIn.m_row[0]); \n "
" out.m_row[1] = fabs(matIn.m_row[1]); \n "
" out.m_row[2] = fabs(matIn.m_row[2]); \n "
" return out; \n "
" } \n "
" __inline \n "
" b3Mat3x3 mtZero(); \n "
" __inline \n "
" b3Mat3x3 mtIdentity(); \n "
" __inline \n "
" b3Mat3x3 mtTranspose(b3Mat3x3 m); \n "
" __inline \n "
" b3Mat3x3 mtMul(b3Mat3x3 a, b3Mat3x3 b); \n "
" __inline \n "
" b3Float4 mtMul1(b3Mat3x3 a, b3Float4 b); \n "
" __inline \n "
" b3Float4 mtMul3(b3Float4 a, b3Mat3x3 b); \n "
" __inline \n "
" b3Mat3x3 mtZero() \n "
" { \n "
" b3Mat3x3 m; \n "
" m.m_row[0] = (b3Float4)(0.f); \n "
" m.m_row[1] = (b3Float4)(0.f); \n "
" m.m_row[2] = (b3Float4)(0.f); \n "
" return m; \n "
" } \n "
" __inline \n "
" b3Mat3x3 mtIdentity() \n "
" { \n "
" b3Mat3x3 m; \n "
" m.m_row[0] = (b3Float4)(1,0,0,0); \n "
" m.m_row[1] = (b3Float4)(0,1,0,0); \n "
" m.m_row[2] = (b3Float4)(0,0,1,0); \n "
" return m; \n "
" } \n "
" __inline \n "
" b3Mat3x3 mtTranspose(b3Mat3x3 m) \n "
" { \n "
" b3Mat3x3 out; \n "
" out.m_row[0] = (b3Float4)(m.m_row[0].x, m.m_row[1].x, m.m_row[2].x, 0.f); \n "
" out.m_row[1] = (b3Float4)(m.m_row[0].y, m.m_row[1].y, m.m_row[2].y, 0.f); \n "
" out.m_row[2] = (b3Float4)(m.m_row[0].z, m.m_row[1].z, m.m_row[2].z, 0.f); \n "
" return out; \n "
" } \n "
" __inline \n "
" b3Mat3x3 mtMul(b3Mat3x3 a, b3Mat3x3 b) \n "
" { \n "
" b3Mat3x3 transB; \n "
" transB = mtTranspose( b ); \n "
" b3Mat3x3 ans; \n "
" // why this doesn't run when 0ing in the for{} \n "
" a.m_row[0].w = 0.f; \n "
" a.m_row[1].w = 0.f; \n "
" a.m_row[2].w = 0.f; \n "
" for(int i=0; i<3; i++) \n "
" { \n "
" // a.m_row[i].w = 0.f; \n "
" ans.m_row[i].x = b3Dot3F4(a.m_row[i],transB.m_row[0]); \n "
" ans.m_row[i].y = b3Dot3F4(a.m_row[i],transB.m_row[1]); \n "
" ans.m_row[i].z = b3Dot3F4(a.m_row[i],transB.m_row[2]); \n "
" ans.m_row[i].w = 0.f; \n "
" } \n "
" return ans; \n "
" } \n "
" __inline \n "
" b3Float4 mtMul1(b3Mat3x3 a, b3Float4 b) \n "
" { \n "
" b3Float4 ans; \n "
" ans.x = b3Dot3F4( a.m_row[0], b ); \n "
" ans.y = b3Dot3F4( a.m_row[1], b ); \n "
" ans.z = b3Dot3F4( a.m_row[2], b ); \n "
" ans.w = 0.f; \n "
" return ans; \n "
" } \n "
" __inline \n "
" b3Float4 mtMul3(b3Float4 a, b3Mat3x3 b) \n "
" { \n "
" b3Float4 colx = b3MakeFloat4(b.m_row[0].x, b.m_row[1].x, b.m_row[2].x, 0); \n "
" b3Float4 coly = b3MakeFloat4(b.m_row[0].y, b.m_row[1].y, b.m_row[2].y, 0); \n "
" b3Float4 colz = b3MakeFloat4(b.m_row[0].z, b.m_row[1].z, b.m_row[2].z, 0); \n "
" b3Float4 ans; \n "
" ans.x = b3Dot3F4( a, colx ); \n "
" ans.y = b3Dot3F4( a, coly ); \n "
" ans.z = b3Dot3F4( a, colz ); \n "
" return ans; \n "
" } \n "
" #endif \n "
" #endif //B3_MAT3x3_H \n "
" typedef struct b3Aabb b3Aabb_t; \n "
" struct b3Aabb \n "
" { \n "
" union \n "
" { \n "
" float m_min[4]; \n "
" b3Float4 m_minVec; \n "
" int m_minIndices[4]; \n "
" }; \n "
" union \n "
" { \n "
" float m_max[4]; \n "
" b3Float4 m_maxVec; \n "
" int m_signedMaxIndices[4]; \n "
" }; \n "
" }; \n "
" inline void b3TransformAabb2(b3Float4ConstArg localAabbMin,b3Float4ConstArg localAabbMax, float margin, \n "
" b3Float4ConstArg pos, \n "
" b3QuatConstArg orn, \n "
" b3Float4* aabbMinOut,b3Float4* aabbMaxOut) \n "
" { \n "
" b3Float4 localHalfExtents = 0.5f*(localAabbMax-localAabbMin); \n "
" localHalfExtents+=b3MakeFloat4(margin,margin,margin,0.f); \n "
" b3Float4 localCenter = 0.5f*(localAabbMax+localAabbMin); \n "
" b3Mat3x3 m; \n "
" m = b3QuatGetRotationMatrix(orn); \n "
" b3Mat3x3 abs_b = b3AbsoluteMat3x3(m); \n "
" b3Float4 center = b3TransformPoint(localCenter,pos,orn); \n "
" \n "
" b3Float4 extent = b3MakeFloat4(b3Dot3F4(localHalfExtents,b3GetRow(abs_b,0)), \n "
" b3Dot3F4(localHalfExtents,b3GetRow(abs_b,1)), \n "
" b3Dot3F4(localHalfExtents,b3GetRow(abs_b,2)), \n "
" 0.f); \n "
" *aabbMinOut = center-extent; \n "
" *aabbMaxOut = center+extent; \n "
" } \n "
" /// conservative test for overlap between two aabbs \n "
" inline bool b3TestAabbAgainstAabb(b3Float4ConstArg aabbMin1,b3Float4ConstArg aabbMax1, \n "
" b3Float4ConstArg aabbMin2, b3Float4ConstArg aabbMax2) \n "
" { \n "
" bool overlap = true; \n "
" overlap = (aabbMin1.x > aabbMax2.x || aabbMax1.x < aabbMin2.x) ? false : overlap; \n "
" overlap = (aabbMin1.z > aabbMax2.z || aabbMax1.z < aabbMin2.z) ? false : overlap; \n "
" overlap = (aabbMin1.y > aabbMax2.y || aabbMax1.y < aabbMin2.y) ? false : overlap; \n "
" return overlap; \n "
" } \n "
" #endif //B3_AABB_H \n "
" /* \n "
" Bullet Continuous Collision Detection and Physics Library \n "
" Copyright (c) 2003-2013 Erwin Coumans http://bulletphysics.org \n "
" This software is provided 'as-is', without any express or implied warranty. \n "
" In no event will the authors be held liable for any damages arising from the use of this software. \n "
" Permission is granted to anyone to use this software for any purpose, \n "
" including commercial applications, and to alter it and redistribute it freely, \n "
" subject to the following restrictions: \n "
" 1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required. \n "
" 2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software. \n "
" 3. This notice may not be removed or altered from any source distribution. \n "
" */ \n "
" #ifndef B3_INT2_H \n "
" #define B3_INT2_H \n "
" #ifdef __cplusplus \n "
" #else \n "
" #define b3UnsignedInt2 uint2 \n "
" #define b3Int2 int2 \n "
" #define b3MakeInt2 (int2) \n "
" #endif //__cplusplus \n "
" #endif \n "
" typedef struct \n "
" { \n "
" float4 m_plane; \n "
" int m_indexOffset; \n "
" int m_numIndices; \n "
" } btGpuFace; \n "
" #define make_float4 (float4) \n "
" __inline \n "
" float4 cross3(float4 a, float4 b) \n "
" { \n "
" return cross(a,b); \n "
" \n "
" // float4 a1 = make_float4(a.xyz,0.f); \n "
" // float4 b1 = make_float4(b.xyz,0.f); \n "
" // return cross(a1,b1); \n "
" //float4 c = make_float4(a.y*b.z - a.z*b.y,a.z*b.x - a.x*b.z,a.x*b.y - a.y*b.x,0.f); \n "
" \n "
" // float4 c = make_float4(a.y*b.z - a.z*b.y,1.f,a.x*b.y - a.y*b.x,0.f); \n "
" \n "
" //return c; \n "
" } \n "
" __inline \n "
" float dot3F4(float4 a, float4 b) \n "
" { \n "
" float4 a1 = make_float4(a.xyz,0.f); \n "
" float4 b1 = make_float4(b.xyz,0.f); \n "
" return dot(a1, b1); \n "
" } \n "
" __inline \n "
" float4 fastNormalize4(float4 v) \n "
" { \n "
" v = make_float4(v.xyz,0.f); \n "
" return fast_normalize(v); \n "
" } \n "
" /////////////////////////////////////// \n "
" // Quaternion \n "
" /////////////////////////////////////// \n "
" typedef float4 Quaternion; \n "
" __inline \n "
" Quaternion qtMul(Quaternion a, Quaternion b); \n "
" __inline \n "
" Quaternion qtNormalize(Quaternion in); \n "
" __inline \n "
" float4 qtRotate(Quaternion q, float4 vec); \n "
" __inline \n "
" Quaternion qtInvert(Quaternion q); \n "
" __inline \n "
" Quaternion qtMul(Quaternion a, Quaternion b) \n "
" { \n "
" Quaternion ans; \n "
" ans = cross3( a, b ); \n "
" ans += a.w*b+b.w*a; \n "
" // ans.w = a.w*b.w - (a.x*b.x+a.y*b.y+a.z*b.z); \n "
" ans.w = a.w*b.w - dot3F4(a, b); \n "
" return ans; \n "
" } \n "
" __inline \n "
" Quaternion qtNormalize(Quaternion in) \n "
" { \n "
" return fastNormalize4(in); \n "
" // in /= length( in ); \n "
" // return in; \n "
" } \n "
" __inline \n "
" float4 qtRotate(Quaternion q, float4 vec) \n "
" { \n "
" Quaternion qInv = qtInvert( q ); \n "
" float4 vcpy = vec; \n "
" vcpy.w = 0.f; \n "
" float4 out = qtMul(qtMul(q,vcpy),qInv); \n "
" return out; \n "
" } \n "
" __inline \n "
" Quaternion qtInvert(Quaternion q) \n "
" { \n "
" return (Quaternion)(-q.xyz, q.w); \n "
" } \n "
" __inline \n "
" float4 qtInvRotate(const Quaternion q, float4 vec) \n "
" { \n "
" return qtRotate( qtInvert( q ), vec ); \n "
" } \n "
" __inline \n "
" float4 transform(const float4* p, const float4* translation, const Quaternion* orientation) \n "
" { \n "
" return qtRotate( *orientation, *p ) + (*translation); \n "
" } \n "
" __inline \n "
" float4 normalize3(const float4 a) \n "
" { \n "
" float4 n = make_float4(a.x, a.y, a.z, 0.f); \n "
" return fastNormalize4( n ); \n "
" } \n "
" inline void projectLocal(const ConvexPolyhedronCL* hull, const float4 pos, const float4 orn, \n "
" const float4* dir, const float4* vertices, float* min, float* max) \n "
" { \n "
" min[0] = FLT_MAX; \n "
" max[0] = -FLT_MAX; \n "
" int numVerts = hull->m_numVertices; \n "
" const float4 localDir = qtInvRotate(orn,*dir); \n "
" float offset = dot(pos,*dir); \n "
" for(int i=0;i<numVerts;i++) \n "
" { \n "
" float dp = dot(vertices[hull->m_vertexOffset+i],localDir); \n "
" if(dp < min[0]) \n "
" min[0] = dp; \n "
" if(dp > max[0]) \n "
" max[0] = dp; \n "
" } \n "
" if(min[0]>max[0]) \n "
" { \n "
" float tmp = min[0]; \n "
" min[0] = max[0]; \n "
" max[0] = tmp; \n "
" } \n "
" min[0] += offset; \n "
" max[0] += offset; \n "
" } \n "
" inline void project(__global const ConvexPolyhedronCL* hull, const float4 pos, const float4 orn, \n "
" const float4* dir, __global const float4* vertices, float* min, float* max) \n "
" { \n "
" min[0] = FLT_MAX; \n "
" max[0] = -FLT_MAX; \n "
" int numVerts = hull->m_numVertices; \n "
" const float4 localDir = qtInvRotate(orn,*dir); \n "
" float offset = dot(pos,*dir); \n "
" for(int i=0;i<numVerts;i++) \n "
" { \n "
" float dp = dot(vertices[hull->m_vertexOffset+i],localDir); \n "
" if(dp < min[0]) \n "
" min[0] = dp; \n "
" if(dp > max[0]) \n "
" max[0] = dp; \n "
" } \n "
" if(min[0]>max[0]) \n "
" { \n "
" float tmp = min[0]; \n "
" min[0] = max[0]; \n "
" max[0] = tmp; \n "
" } \n "
" min[0] += offset; \n "
" max[0] += offset; \n "
" } \n "
" inline bool TestSepAxisLocalA(const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n "
" const float4 posA,const float4 ornA, \n "
" const float4 posB,const float4 ornB, \n "
" float4* sep_axis, const float4* verticesA, __global const float4* verticesB,float* depth) \n "
" { \n "
" float Min0,Max0; \n "
" float Min1,Max1; \n "
" projectLocal(hullA,posA,ornA,sep_axis,verticesA, &Min0, &Max0); \n "
" project(hullB,posB,ornB, sep_axis,verticesB, &Min1, &Max1); \n "
" if(Max0<Min1 || Max1<Min0) \n "
" return false; \n "
" float d0 = Max0 - Min1; \n "
" float d1 = Max1 - Min0; \n "
" *depth = d0<d1 ? d0:d1; \n "
" return true; \n "
" } \n "
" inline bool IsAlmostZero(const float4 v) \n "
" { \n "
" if(fabs(v.x)>1e-6f || fabs(v.y)>1e-6f || fabs(v.z)>1e-6f) \n "
" return false; \n "
" return true; \n "
" } \n "
" bool findSeparatingAxisLocalA( const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n "
" const float4 posA1, \n "
" const float4 ornA, \n "
" const float4 posB1, \n "
" const float4 ornB, \n "
" const float4 DeltaC2, \n "
" \n "
" const float4* verticesA, \n "
" const float4* uniqueEdgesA, \n "
" const btGpuFace* facesA, \n "
" const int* indicesA, \n "
" __global const float4* verticesB, \n "
" __global const float4* uniqueEdgesB, \n "
" __global const btGpuFace* facesB, \n "
" __global const int* indicesB, \n "
" float4* sep, \n "
" float* dmin) \n "
" { \n "
" \n "
" float4 posA = posA1; \n "
" posA.w = 0.f; \n "
" float4 posB = posB1; \n "
" posB.w = 0.f; \n "
" int curPlaneTests=0; \n "
" { \n "
" int numFacesA = hullA->m_numFaces; \n "
" // Test normals from hullA \n "
" for(int i=0;i<numFacesA;i++) \n "
" { \n "
" const float4 normal = facesA[hullA->m_faceOffset+i].m_plane; \n "
" float4 faceANormalWS = qtRotate(ornA,normal); \n "
" if (dot3F4(DeltaC2,faceANormalWS)<0) \n "
" faceANormalWS*=-1.f; \n "
" curPlaneTests++; \n "
" float d; \n "
" if(!TestSepAxisLocalA( hullA, hullB, posA,ornA,posB,ornB,&faceANormalWS, verticesA, verticesB,&d)) \n "
" return false; \n "
" if(d<*dmin) \n "
" { \n "
" *dmin = d; \n "
" *sep = faceANormalWS; \n "
" } \n "
" } \n "
" } \n "
" if((dot3F4(-DeltaC2,*sep))>0.0f) \n "
" { \n "
" *sep = -(*sep); \n "
" } \n "
" return true; \n "
" } \n "
" bool findSeparatingAxisLocalB( __global const ConvexPolyhedronCL* hullA, const ConvexPolyhedronCL* hullB, \n "
" const float4 posA1, \n "
" const float4 ornA, \n "
" const float4 posB1, \n "
" const float4 ornB, \n "
" const float4 DeltaC2, \n "
" __global const float4* verticesA, \n "
" __global const float4* uniqueEdgesA, \n "
" __global const btGpuFace* facesA, \n "
" __global const int* indicesA, \n "
" const float4* verticesB, \n "
" const float4* uniqueEdgesB, \n "
" const btGpuFace* facesB, \n "
" const int* indicesB, \n "
" float4* sep, \n "
" float* dmin) \n "
" { \n "
" float4 posA = posA1; \n "
" posA.w = 0.f; \n "
" float4 posB = posB1; \n "
" posB.w = 0.f; \n "
" int curPlaneTests=0; \n "
" { \n "
" int numFacesA = hullA->m_numFaces; \n "
" // Test normals from hullA \n "
" for(int i=0;i<numFacesA;i++) \n "
" { \n "
" const float4 normal = facesA[hullA->m_faceOffset+i].m_plane; \n "
" float4 faceANormalWS = qtRotate(ornA,normal); \n "
" if (dot3F4(DeltaC2,faceANormalWS)<0) \n "
" faceANormalWS *= -1.f; \n "
" curPlaneTests++; \n "
" float d; \n "
" if(!TestSepAxisLocalA( hullB, hullA, posB,ornB,posA,ornA, &faceANormalWS, verticesB,verticesA, &d)) \n "
" return false; \n "
" if(d<*dmin) \n "
" { \n "
" *dmin = d; \n "
" *sep = faceANormalWS; \n "
" } \n "
" } \n "
" } \n "
" if((dot3F4(-DeltaC2,*sep))>0.0f) \n "
" { \n "
" *sep = -(*sep); \n "
" } \n "
" return true; \n "
" } \n "
" bool findSeparatingAxisEdgeEdgeLocalA( const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n "
" const float4 posA1, \n "
" const float4 ornA, \n "
" const float4 posB1, \n "
" const float4 ornB, \n "
" const float4 DeltaC2, \n "
" const float4* verticesA, \n "
" const float4* uniqueEdgesA, \n "
" const btGpuFace* facesA, \n "
" const int* indicesA, \n "
" __global const float4* verticesB, \n "
" __global const float4* uniqueEdgesB, \n "
" __global const btGpuFace* facesB, \n "
" __global const int* indicesB, \n "
" float4* sep, \n "
" float* dmin) \n "
" { \n "
" float4 posA = posA1; \n "
" posA.w = 0.f; \n "
" float4 posB = posB1; \n "
" posB.w = 0.f; \n "
" int curPlaneTests=0; \n "
" int curEdgeEdge = 0; \n "
" // Test edges \n "
" for(int e0=0;e0<hullA->m_numUniqueEdges;e0++) \n "
" { \n "
" const float4 edge0 = uniqueEdgesA[hullA->m_uniqueEdgesOffset+e0]; \n "
" float4 edge0World = qtRotate(ornA,edge0); \n "
" for(int e1=0;e1<hullB->m_numUniqueEdges;e1++) \n "
" { \n "
" const float4 edge1 = uniqueEdgesB[hullB->m_uniqueEdgesOffset+e1]; \n "
" float4 edge1World = qtRotate(ornB,edge1); \n "
" float4 crossje = cross3(edge0World,edge1World); \n "
" curEdgeEdge++; \n "
" if(!IsAlmostZero(crossje)) \n "
" { \n "
" crossje = normalize3(crossje); \n "
" if (dot3F4(DeltaC2,crossje)<0) \n "
" crossje *= -1.f; \n "
" float dist; \n "
" bool result = true; \n "
" { \n "
" float Min0,Max0; \n "
" float Min1,Max1; \n "
" projectLocal(hullA,posA,ornA,&crossje,verticesA, &Min0, &Max0); \n "
" project(hullB,posB,ornB,&crossje,verticesB, &Min1, &Max1); \n "
" \n "
" if(Max0<Min1 || Max1<Min0) \n "
" result = false; \n "
" \n "
" float d0 = Max0 - Min1; \n "
" float d1 = Max1 - Min0; \n "
" dist = d0<d1 ? d0:d1; \n "
" result = true; \n "
" } \n "
" \n "
" if(dist<*dmin) \n "
" { \n "
" *dmin = dist; \n "
" *sep = crossje; \n "
" } \n "
" } \n "
" } \n "
" } \n "
" \n "
" if((dot3F4(-DeltaC2,*sep))>0.0f) \n "
" { \n "
" *sep = -(*sep); \n "
" } \n "
" return true; \n "
" } \n "
" inline bool TestSepAxis(__global const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n "
" const float4 posA,const float4 ornA, \n "
" const float4 posB,const float4 ornB, \n "
" float4* sep_axis, __global const float4* vertices,float* depth) \n "
" { \n "
" float Min0,Max0; \n "
" float Min1,Max1; \n "
" project(hullA,posA,ornA,sep_axis,vertices, &Min0, &Max0); \n "
" project(hullB,posB,ornB, sep_axis,vertices, &Min1, &Max1); \n "
" if(Max0<Min1 || Max1<Min0) \n "
" return false; \n "
" float d0 = Max0 - Min1; \n "
" float d1 = Max1 - Min0; \n "
" *depth = d0<d1 ? d0:d1; \n "
" return true; \n "
" } \n "
" bool findSeparatingAxis( __global const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n "
" const float4 posA1, \n "
" const float4 ornA, \n "
" const float4 posB1, \n "
" const float4 ornB, \n "
" const float4 DeltaC2, \n "
" __global const float4* vertices, \n "
" __global const float4* uniqueEdges, \n "
" __global const btGpuFace* faces, \n "
" __global const int* indices, \n "
" float4* sep, \n "
" float* dmin) \n "
" { \n "
" \n "
" float4 posA = posA1; \n "
" posA.w = 0.f; \n "
" float4 posB = posB1; \n "
" posB.w = 0.f; \n "
" \n "
" int curPlaneTests=0; \n "
" { \n "
" int numFacesA = hullA->m_numFaces; \n "
" // Test normals from hullA \n "
" for(int i=0;i<numFacesA;i++) \n "
" { \n "
" const float4 normal = faces[hullA->m_faceOffset+i].m_plane; \n "
" float4 faceANormalWS = qtRotate(ornA,normal); \n "
" \n "
" if (dot3F4(DeltaC2,faceANormalWS)<0) \n "
" faceANormalWS*=-1.f; \n "
" \n "
" curPlaneTests++; \n "
" \n "
" float d; \n "
" if(!TestSepAxis( hullA, hullB, posA,ornA,posB,ornB,&faceANormalWS, vertices,&d)) \n "
" return false; \n "
" \n "
" if(d<*dmin) \n "
" { \n "
" *dmin = d; \n "
" *sep = faceANormalWS; \n "
" } \n "
" } \n "
" } \n "
" if((dot3F4(-DeltaC2,*sep))>0.0f) \n "
" { \n "
" *sep = -(*sep); \n "
" } \n "
" \n "
" return true; \n "
" } \n "
" bool findSeparatingAxisUnitSphere( __global const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n "
" const float4 posA1, \n "
" const float4 ornA, \n "
" const float4 posB1, \n "
" const float4 ornB, \n "
" const float4 DeltaC2, \n "
" __global const float4* vertices, \n "
" __global const float4* unitSphereDirections, \n "
" int numUnitSphereDirections, \n "
" float4* sep, \n "
" float* dmin) \n "
" { \n "
" \n "
" float4 posA = posA1; \n "
" posA.w = 0.f; \n "
" float4 posB = posB1; \n "
" posB.w = 0.f; \n "
" int curPlaneTests=0; \n "
" int curEdgeEdge = 0; \n "
" // Test unit sphere directions \n "
" for (int i=0;i<numUnitSphereDirections;i++) \n "
" { \n "
" float4 crossje; \n "
" crossje = unitSphereDirections[i]; \n "
" if (dot3F4(DeltaC2,crossje)>0) \n "
" crossje *= -1.f; \n "
" { \n "
" float dist; \n "
" bool result = true; \n "
" float Min0,Max0; \n "
" float Min1,Max1; \n "
" project(hullA,posA,ornA,&crossje,vertices, &Min0, &Max0); \n "
" project(hullB,posB,ornB,&crossje,vertices, &Min1, &Max1); \n "
" \n "
" if(Max0<Min1 || Max1<Min0) \n "
" return false; \n "
" \n "
" float d0 = Max0 - Min1; \n "
" float d1 = Max1 - Min0; \n "
" dist = d0<d1 ? d0:d1; \n "
" result = true; \n "
" \n "
" if(dist<*dmin) \n "
" { \n "
" *dmin = dist; \n "
" *sep = crossje; \n "
" } \n "
" } \n "
" } \n "
" \n "
" if((dot3F4(-DeltaC2,*sep))>0.0f) \n "
" { \n "
" *sep = -(*sep); \n "
" } \n "
" return true; \n "
" } \n "
" bool findSeparatingAxisEdgeEdge( __global const ConvexPolyhedronCL* hullA, __global const ConvexPolyhedronCL* hullB, \n "
" const float4 posA1, \n "
" const float4 ornA, \n "
" const float4 posB1, \n "
" const float4 ornB, \n "
" const float4 DeltaC2, \n "
" __global const float4* vertices, \n "
" __global const float4* uniqueEdges, \n "
" __global const btGpuFace* faces, \n "
" __global const int* indices, \n "
" float4* sep, \n "
" float* dmin) \n "
" { \n "
" \n "
" float4 posA = posA1; \n "
" posA.w = 0.f; \n "
" float4 posB = posB1; \n "
" posB.w = 0.f; \n "
" int curPlaneTests=0; \n "
" int curEdgeEdge = 0; \n "
" // Test edges \n "
" for(int e0=0;e0<hullA->m_numUniqueEdges;e0++) \n "
" { \n "
" const float4 edge0 = uniqueEdges[hullA->m_uniqueEdgesOffset+e0]; \n "
" float4 edge0World = qtRotate(ornA,edge0); \n "
" for(int e1=0;e1<hullB->m_numUniqueEdges;e1++) \n "
" { \n "
" const float4 edge1 = uniqueEdges[hullB->m_uniqueEdgesOffset+e1]; \n "
" float4 edge1World = qtRotate(ornB,edge1); \n "
" float4 crossje = cross3(edge0World,edge1World); \n "
" curEdgeEdge++; \n "
" if(!IsAlmostZero(crossje)) \n "
" { \n "
" crossje = normalize3(crossje); \n "
" if (dot3F4(DeltaC2,crossje)<0) \n "
" crossje*=-1.f; \n "
" \n "
" float dist; \n "
" bool result = true; \n "
" { \n "
" float Min0,Max0; \n "
" float Min1,Max1; \n "
" project(hullA,posA,ornA,&crossje,vertices, &Min0, &Max0); \n "
" project(hullB,posB,ornB,&crossje,vertices, &Min1, &Max1); \n "
" \n "
" if(Max0<Min1 || Max1<Min0) \n "
" return false; \n "
" \n "
" float d0 = Max0 - Min1; \n "
" float d1 = Max1 - Min0; \n "
" dist = d0<d1 ? d0:d1; \n "
" result = true; \n "
" } \n "
" \n "
" if(dist<*dmin) \n "
" { \n "
" *dmin = dist; \n "
" *sep = crossje; \n "
" } \n "
" } \n "
" } \n "
" } \n "
" \n "
" if((dot3F4(-DeltaC2,*sep))>0.0f) \n "
" { \n "
" *sep = -(*sep); \n "
" } \n "
" return true; \n "
" } \n "
" // work-in-progress \n "
" __kernel void processCompoundPairsKernel( __global const int4* gpuCompoundPairs, \n "
" __global const BodyData* rigidBodies, \n "
" __global const btCollidableGpu* collidables, \n "
" __global const ConvexPolyhedronCL* convexShapes, \n "
" __global const float4* vertices, \n "
" __global const float4* uniqueEdges, \n "
" __global const btGpuFace* faces, \n "
" __global const int* indices, \n "
" __global btAabbCL* aabbs, \n "
" __global const btGpuChildShape* gpuChildShapes, \n "
" __global volatile float4* gpuCompoundSepNormalsOut, \n "
" __global volatile int* gpuHasCompoundSepNormalsOut, \n "
" int numCompoundPairs \n "
" ) \n "
" { \n "
" int i = get_global_id(0); \n "
" if (i<numCompoundPairs) \n "
" { \n "
" int bodyIndexA = gpuCompoundPairs[i].x; \n "
" int bodyIndexB = gpuCompoundPairs[i].y; \n "
" int childShapeIndexA = gpuCompoundPairs[i].z; \n "
" int childShapeIndexB = gpuCompoundPairs[i].w; \n "
" \n "
" int collidableIndexA = -1; \n "
" int collidableIndexB = -1; \n "
" \n "
" float4 ornA = rigidBodies[bodyIndexA].m_quat; \n "
" float4 posA = rigidBodies[bodyIndexA].m_pos; \n "
" \n "
" float4 ornB = rigidBodies[bodyIndexB].m_quat; \n "
" float4 posB = rigidBodies[bodyIndexB].m_pos; \n "
" \n "
" if (childShapeIndexA >= 0) \n "
" { \n "
" collidableIndexA = gpuChildShapes[childShapeIndexA].m_shapeIndex; \n "
" float4 childPosA = gpuChildShapes[childShapeIndexA].m_childPosition; \n "
" float4 childOrnA = gpuChildShapes[childShapeIndexA].m_childOrientation; \n "
" float4 newPosA = qtRotate(ornA,childPosA)+posA; \n "
" float4 newOrnA = qtMul(ornA,childOrnA); \n "
" posA = newPosA; \n "
" ornA = newOrnA; \n "
" } else \n "
" { \n "
" collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx; \n "
" } \n "
" \n "
" if (childShapeIndexB>=0) \n "
" { \n "
" collidableIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex; \n "
" float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition; \n "
" float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation; \n "
" float4 newPosB = transform(&childPosB,&posB,&ornB); \n "
" float4 newOrnB = qtMul(ornB,childOrnB); \n "
" posB = newPosB; \n "
" ornB = newOrnB; \n "
" } else \n "
" { \n "
" collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx; \n "
" } \n "
" \n "
" gpuHasCompoundSepNormalsOut[i] = 0; \n "
" \n "
" int shapeIndexA = collidables[collidableIndexA].m_shapeIndex; \n "
" int shapeIndexB = collidables[collidableIndexB].m_shapeIndex; \n "
" \n "
" int shapeTypeA = collidables[collidableIndexA].m_shapeType; \n "
" int shapeTypeB = collidables[collidableIndexB].m_shapeType; \n "
" \n "
" if ((shapeTypeA != SHAPE_CONVEX_HULL) || (shapeTypeB != SHAPE_CONVEX_HULL)) \n "
" { \n "
" return; \n "
" } \n "
" int hasSeparatingAxis = 5; \n "
" \n "
" int numFacesA = convexShapes[shapeIndexA].m_numFaces; \n "
" float dmin = FLT_MAX; \n "
" posA.w = 0.f; \n "
" posB.w = 0.f; \n "
" float4 c0local = convexShapes[shapeIndexA].m_localCenter; \n "
" float4 c0 = transform(&c0local, &posA, &ornA); \n "
" float4 c1local = convexShapes[shapeIndexB].m_localCenter; \n "
" float4 c1 = transform(&c1local,&posB,&ornB); \n "
" const float4 DeltaC2 = c0 - c1; \n "
" float4 sepNormal = make_float4(1,0,0,0); \n "
" bool sepA = findSeparatingAxis( &convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA,posB,ornB,DeltaC2,vertices,uniqueEdges,faces,indices,&sepNormal,&dmin); \n "
" hasSeparatingAxis = 4; \n "
" if (!sepA) \n "
" { \n "
" hasSeparatingAxis = 0; \n "
" } else \n "
" { \n "
" bool sepB = findSeparatingAxis( &convexShapes[shapeIndexB],&convexShapes[shapeIndexA],posB,ornB,posA,ornA,DeltaC2,vertices,uniqueEdges,faces,indices,&sepNormal,&dmin); \n "
" if (!sepB) \n "
" { \n "
" hasSeparatingAxis = 0; \n "
" } else//(!sepB) \n "
" { \n "
" bool sepEE = findSeparatingAxisEdgeEdge( &convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA,posB,ornB,DeltaC2,vertices,uniqueEdges,faces,indices,&sepNormal,&dmin); \n "
" if (sepEE) \n "
" { \n "
" gpuCompoundSepNormalsOut[i] = sepNormal;//fastNormalize4(sepNormal); \n "
" gpuHasCompoundSepNormalsOut[i] = 1; \n "
" }//sepEE \n "
" }//(!sepB) \n "
" }//(!sepA) \n "
" \n "
" \n "
" } \n "
" \n "
" } \n "
" inline b3Float4 MyUnQuantize(const unsigned short* vecIn, b3Float4 quantization, b3Float4 bvhAabbMin) \n "
" { \n "
" b3Float4 vecOut; \n "
" vecOut = b3MakeFloat4( \n "
" (float)(vecIn[0]) / (quantization.x), \n "
" (float)(vecIn[1]) / (quantization.y), \n "
" (float)(vecIn[2]) / (quantization.z), \n "
" 0.f); \n "
" vecOut += bvhAabbMin; \n "
" return vecOut; \n "
" } \n "
" inline b3Float4 MyUnQuantizeGlobal(__global const unsigned short* vecIn, b3Float4 quantization, b3Float4 bvhAabbMin) \n "
" { \n "
" b3Float4 vecOut; \n "
" vecOut = b3MakeFloat4( \n "
" (float)(vecIn[0]) / (quantization.x), \n "
" (float)(vecIn[1]) / (quantization.y), \n "
" (float)(vecIn[2]) / (quantization.z), \n "
" 0.f); \n "
" vecOut += bvhAabbMin; \n "
" return vecOut; \n "
" } \n "
" // work-in-progress \n "
" __kernel void findCompoundPairsKernel( __global const int4* pairs, \n "
" __global const BodyData* rigidBodies, \n "
" __global const btCollidableGpu* collidables, \n "
" __global const ConvexPolyhedronCL* convexShapes, \n "
" __global const float4* vertices, \n "
" __global const float4* uniqueEdges, \n "
" __global const btGpuFace* faces, \n "
" __global const int* indices, \n "
" __global b3Aabb_t* aabbLocalSpace, \n "
" __global const btGpuChildShape* gpuChildShapes, \n "
" __global volatile int4* gpuCompoundPairsOut, \n "
" __global volatile int* numCompoundPairsOut, \n "
" __global const b3BvhSubtreeInfo* subtrees, \n "
" __global const b3QuantizedBvhNode* quantizedNodes, \n "
" __global const b3BvhInfo* bvhInfos, \n "
" int numPairs, \n "
" int maxNumCompoundPairsCapacity \n "
" ) \n "
" { \n "
" int i = get_global_id(0); \n "
" if (i<numPairs) \n "
" { \n "
" int bodyIndexA = pairs[i].x; \n "
" int bodyIndexB = pairs[i].y; \n "
" int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx; \n "
" int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx; \n "
" int shapeIndexA = collidables[collidableIndexA].m_shapeIndex; \n "
" int shapeIndexB = collidables[collidableIndexB].m_shapeIndex; \n "
" //once the broadphase avoids static-static pairs, we can remove this test \n "
" if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0)) \n "
" { \n "
" return; \n "
" } \n "
" if ((collidables[collidableIndexA].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) &&(collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)) \n "
" { \n "
" int bvhA = collidables[collidableIndexA].m_compoundBvhIndex; \n "
" int bvhB = collidables[collidableIndexB].m_compoundBvhIndex; \n "
" int numSubTreesA = bvhInfos[bvhA].m_numSubTrees; \n "
" int subTreesOffsetA = bvhInfos[bvhA].m_subTreeOffset; \n "
" int subTreesOffsetB = bvhInfos[bvhB].m_subTreeOffset; \n "
" int numSubTreesB = bvhInfos[bvhB].m_numSubTrees; \n "
" \n "
" float4 posA = rigidBodies[bodyIndexA].m_pos; \n "
" b3Quat ornA = rigidBodies[bodyIndexA].m_quat; \n "
" b3Quat ornB = rigidBodies[bodyIndexB].m_quat; \n "
" float4 posB = rigidBodies[bodyIndexB].m_pos; \n "
" \n "
" for (int p=0;p<numSubTreesA;p++) \n "
" { \n "
" b3BvhSubtreeInfo subtreeA = subtrees[subTreesOffsetA+p]; \n "
" //bvhInfos[bvhA].m_quantization \n "
" b3Float4 treeAminLocal = MyUnQuantize(subtreeA.m_quantizedAabbMin,bvhInfos[bvhA].m_quantization,bvhInfos[bvhA].m_aabbMin); \n "
" b3Float4 treeAmaxLocal = MyUnQuantize(subtreeA.m_quantizedAabbMax,bvhInfos[bvhA].m_quantization,bvhInfos[bvhA].m_aabbMin); \n "
" b3Float4 aabbAMinOut,aabbAMaxOut; \n "
" float margin=0.f; \n "
" b3TransformAabb2(treeAminLocal,treeAmaxLocal, margin,posA,ornA,&aabbAMinOut,&aabbAMaxOut); \n "
" \n "
" for (int q=0;q<numSubTreesB;q++) \n "
" { \n "
" b3BvhSubtreeInfo subtreeB = subtrees[subTreesOffsetB+q]; \n "
" b3Float4 treeBminLocal = MyUnQuantize(subtreeB.m_quantizedAabbMin,bvhInfos[bvhB].m_quantization,bvhInfos[bvhB].m_aabbMin); \n "
" b3Float4 treeBmaxLocal = MyUnQuantize(subtreeB.m_quantizedAabbMax,bvhInfos[bvhB].m_quantization,bvhInfos[bvhB].m_aabbMin); \n "
" b3Float4 aabbBMinOut,aabbBMaxOut; \n "
" float margin=0.f; \n "
" b3TransformAabb2(treeBminLocal,treeBmaxLocal, margin,posB,ornB,&aabbBMinOut,&aabbBMaxOut); \n "
" \n "
" \n "
" bool aabbOverlap = b3TestAabbAgainstAabb(aabbAMinOut,aabbAMaxOut,aabbBMinOut,aabbBMaxOut); \n "
" if (aabbOverlap) \n "
" { \n "
" \n "
" int startNodeIndexA = subtreeA.m_rootNodeIndex+bvhInfos[bvhA].m_nodeOffset; \n "
" int endNodeIndexA = startNodeIndexA+subtreeA.m_subtreeSize; \n "
" int startNodeIndexB = subtreeB.m_rootNodeIndex+bvhInfos[bvhB].m_nodeOffset; \n "
" int endNodeIndexB = startNodeIndexB+subtreeB.m_subtreeSize; \n "
" b3Int2 nodeStack[B3_MAX_STACK_DEPTH]; \n "
" b3Int2 node0; \n "
" node0.x = startNodeIndexA; \n "
" node0.y = startNodeIndexB; \n "
" int maxStackDepth = B3_MAX_STACK_DEPTH; \n "
" int depth=0; \n "
" nodeStack[depth++]=node0; \n "
" do \n "
" { \n "
" b3Int2 node = nodeStack[--depth]; \n "
" b3Float4 aMinLocal = MyUnQuantizeGlobal(quantizedNodes[node.x].m_quantizedAabbMin,bvhInfos[bvhA].m_quantization,bvhInfos[bvhA].m_aabbMin); \n "
" b3Float4 aMaxLocal = MyUnQuantizeGlobal(quantizedNodes[node.x].m_quantizedAabbMax,bvhInfos[bvhA].m_quantization,bvhInfos[bvhA].m_aabbMin); \n "
" b3Float4 bMinLocal = MyUnQuantizeGlobal(quantizedNodes[node.y].m_quantizedAabbMin,bvhInfos[bvhB].m_quantization,bvhInfos[bvhB].m_aabbMin); \n "
" b3Float4 bMaxLocal = MyUnQuantizeGlobal(quantizedNodes[node.y].m_quantizedAabbMax,bvhInfos[bvhB].m_quantization,bvhInfos[bvhB].m_aabbMin); \n "
" float margin=0.f; \n "
" b3Float4 aabbAMinOut,aabbAMaxOut; \n "
" b3TransformAabb2(aMinLocal,aMaxLocal, margin,posA,ornA,&aabbAMinOut,&aabbAMaxOut); \n "
" b3Float4 aabbBMinOut,aabbBMaxOut; \n "
" b3TransformAabb2(bMinLocal,bMaxLocal, margin,posB,ornB,&aabbBMinOut,&aabbBMaxOut); \n "
" \n "
" bool nodeOverlap = b3TestAabbAgainstAabb(aabbAMinOut,aabbAMaxOut,aabbBMinOut,aabbBMaxOut); \n "
" if (nodeOverlap) \n "
" { \n "
" bool isLeafA = isLeafNodeGlobal(&quantizedNodes[node.x]); \n "
" bool isLeafB = isLeafNodeGlobal(&quantizedNodes[node.y]); \n "
" bool isInternalA = !isLeafA; \n "
" bool isInternalB = !isLeafB; \n "
" //fail, even though it might hit two leaf nodes \n "
" if (depth+4>maxStackDepth && !(isLeafA && isLeafB)) \n "
" { \n "
" //printf( \" Error: traversal exceeded maxStackDepth \" ); \n "
" continue; \n "
" } \n "
" if(isInternalA) \n "
" { \n "
" int nodeAleftChild = node.x+1; \n "
" bool isNodeALeftChildLeaf = isLeafNodeGlobal(&quantizedNodes[node.x+1]); \n "
" int nodeArightChild = isNodeALeftChildLeaf? node.x+2 : node.x+1 + getEscapeIndexGlobal(&quantizedNodes[node.x+1]); \n "
" if(isInternalB) \n "
" { \n "
" int nodeBleftChild = node.y+1; \n "
" bool isNodeBLeftChildLeaf = isLeafNodeGlobal(&quantizedNodes[node.y+1]); \n "
" int nodeBrightChild = isNodeBLeftChildLeaf? node.y+2 : node.y+1 + getEscapeIndexGlobal(&quantizedNodes[node.y+1]); \n "
" nodeStack[depth++] = b3MakeInt2(nodeAleftChild, nodeBleftChild); \n "
" nodeStack[depth++] = b3MakeInt2(nodeArightChild, nodeBleftChild); \n "
" nodeStack[depth++] = b3MakeInt2(nodeAleftChild, nodeBrightChild); \n "
" nodeStack[depth++] = b3MakeInt2(nodeArightChild, nodeBrightChild); \n "
" } \n "
" else \n "
" { \n "
" nodeStack[depth++] = b3MakeInt2(nodeAleftChild,node.y); \n "
" nodeStack[depth++] = b3MakeInt2(nodeArightChild,node.y); \n "
" } \n "
" } \n "
" else \n "
" { \n "
" if(isInternalB) \n "
" { \n "
" int nodeBleftChild = node.y+1; \n "
" bool isNodeBLeftChildLeaf = isLeafNodeGlobal(&quantizedNodes[node.y+1]); \n "
" int nodeBrightChild = isNodeBLeftChildLeaf? node.y+2 : node.y+1 + getEscapeIndexGlobal(&quantizedNodes[node.y+1]); \n "
" nodeStack[depth++] = b3MakeInt2(node.x,nodeBleftChild); \n "
" nodeStack[depth++] = b3MakeInt2(node.x,nodeBrightChild); \n "
" } \n "
" else \n "
" { \n "
" int compoundPairIdx = atomic_inc(numCompoundPairsOut); \n "
" if (compoundPairIdx<maxNumCompoundPairsCapacity) \n "
" { \n "
" int childShapeIndexA = getTriangleIndexGlobal(&quantizedNodes[node.x]); \n "
" int childShapeIndexB = getTriangleIndexGlobal(&quantizedNodes[node.y]); \n "
" gpuCompoundPairsOut[compoundPairIdx] = (int4)(bodyIndexA,bodyIndexB,childShapeIndexA,childShapeIndexB); \n "
" } \n "
" } \n "
" } \n "
" } \n "
" } while (depth); \n "
" } \n "
" } \n "
" } \n "
" \n "
" return; \n "
" } \n "
" if ((collidables[collidableIndexA].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) ||(collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)) \n "
" { \n "
" if (collidables[collidableIndexA].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) \n "
" { \n "
" int numChildrenA = collidables[collidableIndexA].m_numChildShapes; \n "
" for (int c=0;c<numChildrenA;c++) \n "
" { \n "
" int childShapeIndexA = collidables[collidableIndexA].m_shapeIndex+c; \n "
" int childColIndexA = gpuChildShapes[childShapeIndexA].m_shapeIndex; \n "
" float4 posA = rigidBodies[bodyIndexA].m_pos; \n "
" float4 ornA = rigidBodies[bodyIndexA].m_quat; \n "
" float4 childPosA = gpuChildShapes[childShapeIndexA].m_childPosition; \n "
" float4 childOrnA = gpuChildShapes[childShapeIndexA].m_childOrientation; \n "
" float4 newPosA = qtRotate(ornA,childPosA)+posA; \n "
" float4 newOrnA = qtMul(ornA,childOrnA); \n "
" int shapeIndexA = collidables[childColIndexA].m_shapeIndex; \n "
" b3Aabb_t aabbAlocal = aabbLocalSpace[shapeIndexA]; \n "
" float margin = 0.f; \n "
" \n "
" b3Float4 aabbAMinWS; \n "
" b3Float4 aabbAMaxWS; \n "
" \n "
" b3TransformAabb2(aabbAlocal.m_minVec,aabbAlocal.m_maxVec,margin, \n "
" newPosA, \n "
" newOrnA, \n "
" &aabbAMinWS,&aabbAMaxWS); \n "
" \n "
" \n "
" if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) \n "
" { \n "
" int numChildrenB = collidables[collidableIndexB].m_numChildShapes; \n "
" for (int b=0;b<numChildrenB;b++) \n "
" { \n "
" int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+b; \n "
" int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex; \n "
" float4 ornB = rigidBodies[bodyIndexB].m_quat; \n "
" float4 posB = rigidBodies[bodyIndexB].m_pos; \n "
" float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition; \n "
" float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation; \n "
" float4 newPosB = transform(&childPosB,&posB,&ornB); \n "
" float4 newOrnB = qtMul(ornB,childOrnB); \n "
" int shapeIndexB = collidables[childColIndexB].m_shapeIndex; \n "
" b3Aabb_t aabbBlocal = aabbLocalSpace[shapeIndexB]; \n "
" \n "
" b3Float4 aabbBMinWS; \n "
" b3Float4 aabbBMaxWS; \n "
" \n "
" b3TransformAabb2(aabbBlocal.m_minVec,aabbBlocal.m_maxVec,margin, \n "
" newPosB, \n "
" newOrnB, \n "
" &aabbBMinWS,&aabbBMaxWS); \n "
" \n "
" \n "
" \n "
" bool aabbOverlap = b3TestAabbAgainstAabb(aabbAMinWS,aabbAMaxWS,aabbBMinWS,aabbBMaxWS); \n "
" if (aabbOverlap) \n "
" { \n "
" int numFacesA = convexShapes[shapeIndexA].m_numFaces; \n "
" float dmin = FLT_MAX; \n "
" float4 posA = newPosA; \n "
" posA.w = 0.f; \n "
" float4 posB = newPosB; \n "
" posB.w = 0.f; \n "
" float4 c0local = convexShapes[shapeIndexA].m_localCenter; \n "
" float4 ornA = newOrnA; \n "
" float4 c0 = transform(&c0local, &posA, &ornA); \n "
" float4 c1local = convexShapes[shapeIndexB].m_localCenter; \n "
" float4 ornB =newOrnB; \n "
" float4 c1 = transform(&c1local,&posB,&ornB); \n "
" const float4 DeltaC2 = c0 - c1; \n "
" {// \n "
" int compoundPairIdx = atomic_inc(numCompoundPairsOut); \n "
" if (compoundPairIdx<maxNumCompoundPairsCapacity) \n "
" { \n "
" gpuCompoundPairsOut[compoundPairIdx] = (int4)(bodyIndexA,bodyIndexB,childShapeIndexA,childShapeIndexB); \n "
" } \n "
" }// \n "
" }//fi(1) \n "
" } //for (int b=0 \n "
" }//if (collidables[collidableIndexB]. \n "
" else//if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) \n "
" { \n "
" if (1) \n "
" { \n "
" int numFacesA = convexShapes[shapeIndexA].m_numFaces; \n "
" float dmin = FLT_MAX; \n "
" float4 posA = newPosA; \n "
" posA.w = 0.f; \n "
" float4 posB = rigidBodies[bodyIndexB].m_pos; \n "
" posB.w = 0.f; \n "
" float4 c0local = convexShapes[shapeIndexA].m_localCenter; \n "
" float4 ornA = newOrnA; \n "
" float4 c0 = transform(&c0local, &posA, &ornA); \n "
" float4 c1local = convexShapes[shapeIndexB].m_localCenter; \n "
" float4 ornB = rigidBodies[bodyIndexB].m_quat; \n "
" float4 c1 = transform(&c1local,&posB,&ornB); \n "
" const float4 DeltaC2 = c0 - c1; \n "
" { \n "
" int compoundPairIdx = atomic_inc(numCompoundPairsOut); \n "
" if (compoundPairIdx<maxNumCompoundPairsCapacity) \n "
" { \n "
" gpuCompoundPairsOut[compoundPairIdx] = (int4)(bodyIndexA,bodyIndexB,childShapeIndexA,-1); \n "
" }//if (compoundPairIdx<maxNumCompoundPairsCapacity) \n "
" }// \n "
" }//fi (1) \n "
" }//if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) \n "
" }//for (int b=0;b<numChildrenB;b++) \n "
" return; \n "
" }//if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) \n "
" if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONCAVE_TRIMESH) \n "
" && (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)) \n "
" { \n "
" int numChildrenB = collidables[collidableIndexB].m_numChildShapes; \n "
" for (int b=0;b<numChildrenB;b++) \n "
" { \n "
" int childShapeIndexB = collidables[collidableIndexB].m_shapeIndex+b; \n "
" int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex; \n "
" float4 ornB = rigidBodies[bodyIndexB].m_quat; \n "
" float4 posB = rigidBodies[bodyIndexB].m_pos; \n "
" float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition; \n "
" float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation; \n "
" float4 newPosB = qtRotate(ornB,childPosB)+posB; \n "
" float4 newOrnB = qtMul(ornB,childOrnB); \n "
" int shapeIndexB = collidables[childColIndexB].m_shapeIndex; \n "
" ////////////////////////////////////// \n "
" if (1) \n "
" { \n "
" int numFacesA = convexShapes[shapeIndexA].m_numFaces; \n "
" float dmin = FLT_MAX; \n "
" float4 posA = rigidBodies[bodyIndexA].m_pos; \n "
" posA.w = 0.f; \n "
" float4 posB = newPosB; \n "
" posB.w = 0.f; \n "
" float4 c0local = convexShapes[shapeIndexA].m_localCenter; \n "
" float4 ornA = rigidBodies[bodyIndexA].m_quat; \n "
" float4 c0 = transform(&c0local, &posA, &ornA); \n "
" float4 c1local = convexShapes[shapeIndexB].m_localCenter; \n "
" float4 ornB =newOrnB; \n "
" float4 c1 = transform(&c1local,&posB,&ornB); \n "
" const float4 DeltaC2 = c0 - c1; \n "
" {// \n "
" int compoundPairIdx = atomic_inc(numCompoundPairsOut); \n "
" if (compoundPairIdx<maxNumCompoundPairsCapacity) \n "
" { \n "
" gpuCompoundPairsOut[compoundPairIdx] = (int4)(bodyIndexA,bodyIndexB,-1,childShapeIndexB); \n "
" }//fi (compoundPairIdx<maxNumCompoundPairsCapacity) \n "
" }// \n "
" }//fi (1) \n "
" }//for (int b=0;b<numChildrenB;b++) \n "
" return; \n "
" }//if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) \n "
" return; \n "
" }//fi ((collidables[collidableIndexA].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) ||(collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS)) \n "
" }//i<numPairs \n "
" } \n "
" // work-in-progress \n "
" __kernel void findSeparatingAxisKernel( __global const int4* pairs, \n "
" __global const BodyData* rigidBodies, \n "
" __global const btCollidableGpu* collidables, \n "
" __global const ConvexPolyhedronCL* convexShapes, \n "
" __global const float4* vertices, \n "
" __global const float4* uniqueEdges, \n "
" __global const btGpuFace* faces, \n "
" __global const int* indices, \n "
" __global btAabbCL* aabbs, \n "
" __global volatile float4* separatingNormals, \n "
" __global volatile int* hasSeparatingAxis, \n "
" int numPairs \n "
" ) \n "
" { \n "
" int i = get_global_id(0); \n "
" \n "
" if (i<numPairs) \n "
" { \n "
" \n "
" int bodyIndexA = pairs[i].x; \n "
" int bodyIndexB = pairs[i].y; \n "
" int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx; \n "
" int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx; \n "
" \n "
" int shapeIndexA = collidables[collidableIndexA].m_shapeIndex; \n "
" int shapeIndexB = collidables[collidableIndexB].m_shapeIndex; \n "
" \n "
" \n "
" //once the broadphase avoids static-static pairs, we can remove this test \n "
" if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0)) \n "
" { \n "
" hasSeparatingAxis[i] = 0; \n "
" return; \n "
" } \n "
" \n "
" if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONVEX_HULL) ||(collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL)) \n "
" { \n "
" hasSeparatingAxis[i] = 0; \n "
" return; \n "
" } \n "
" \n "
" if ((collidables[collidableIndexA].m_shapeType==SHAPE_CONCAVE_TRIMESH)) \n "
" { \n "
" hasSeparatingAxis[i] = 0; \n "
" return; \n "
" } \n "
" int numFacesA = convexShapes[shapeIndexA].m_numFaces; \n "
" float dmin = FLT_MAX; \n "
" float4 posA = rigidBodies[bodyIndexA].m_pos; \n "
" posA.w = 0.f; \n "
" float4 posB = rigidBodies[bodyIndexB].m_pos; \n "
" posB.w = 0.f; \n "
" float4 c0local = convexShapes[shapeIndexA].m_localCenter; \n "
" float4 ornA = rigidBodies[bodyIndexA].m_quat; \n "
" float4 c0 = transform(&c0local, &posA, &ornA); \n "
" float4 c1local = convexShapes[shapeIndexB].m_localCenter; \n "
" float4 ornB =rigidBodies[bodyIndexB].m_quat; \n "
" float4 c1 = transform(&c1local,&posB,&ornB); \n "
" const float4 DeltaC2 = c0 - c1; \n "
" float4 sepNormal; \n "
" \n "
" bool sepA = findSeparatingAxis( &convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA, \n "
" posB,ornB, \n "
" DeltaC2, \n "
" vertices,uniqueEdges,faces, \n "
" indices,&sepNormal,&dmin); \n "
" hasSeparatingAxis[i] = 4; \n "
" if (!sepA) \n "
" { \n "
" hasSeparatingAxis[i] = 0; \n "
" } else \n "
" { \n "
" bool sepB = findSeparatingAxis( &convexShapes[shapeIndexB],&convexShapes[shapeIndexA],posB,ornB, \n "
" posA,ornA, \n "
" DeltaC2, \n "
" vertices,uniqueEdges,faces, \n "
" indices,&sepNormal,&dmin); \n "
" if (!sepB) \n "
" { \n "
" hasSeparatingAxis[i] = 0; \n "
" } else \n "
" { \n "
" bool sepEE = findSeparatingAxisEdgeEdge( &convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA, \n "
" posB,ornB, \n "
" DeltaC2, \n "
" vertices,uniqueEdges,faces, \n "
" indices,&sepNormal,&dmin); \n "
" if (!sepEE) \n "
" { \n "
" hasSeparatingAxis[i] = 0; \n "
" } else \n "
" { \n "
" hasSeparatingAxis[i] = 1; \n "
" separatingNormals[i] = sepNormal; \n "
" } \n "
" } \n "
" } \n "
" \n "
" } \n "
" } \n "
" __kernel void findSeparatingAxisVertexFaceKernel( __global const int4* pairs, \n "
" __global const BodyData* rigidBodies, \n "
" __global const btCollidableGpu* collidables, \n "
" __global const ConvexPolyhedronCL* convexShapes, \n "
" __global const float4* vertices, \n "
" __global const float4* uniqueEdges, \n "
" __global const btGpuFace* faces, \n "
" __global const int* indices, \n "
" __global btAabbCL* aabbs, \n "
" __global volatile float4* separatingNormals, \n "
" __global volatile int* hasSeparatingAxis, \n "
" __global float* dmins, \n "
" int numPairs \n "
" ) \n "
" { \n "
" int i = get_global_id(0); \n "
" \n "
" if (i<numPairs) \n "
" { \n "
" \n "
" int bodyIndexA = pairs[i].x; \n "
" int bodyIndexB = pairs[i].y; \n "
" int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx; \n "
" int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx; \n "
" \n "
" int shapeIndexA = collidables[collidableIndexA].m_shapeIndex; \n "
" int shapeIndexB = collidables[collidableIndexB].m_shapeIndex; \n "
" \n "
" hasSeparatingAxis[i] = 0; \n "
" \n "
" //once the broadphase avoids static-static pairs, we can remove this test \n "
" if ((rigidBodies[bodyIndexA].m_invMass==0) &&(rigidBodies[bodyIndexB].m_invMass==0)) \n "
" { \n "
" return; \n "
" } \n "
" \n "
" if ((collidables[collidableIndexA].m_shapeType!=SHAPE_CONVEX_HULL) ||(collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL)) \n "
" { \n "
" return; \n "
" } \n "
" \n "
" int numFacesA = convexShapes[shapeIndexA].m_numFaces; \n "
" float dmin = FLT_MAX; \n "
" dmins[i] = dmin; \n "
" \n "
" float4 posA = rigidBodies[bodyIndexA].m_pos; \n "
" posA.w = 0.f; \n "
" float4 posB = rigidBodies[bodyIndexB].m_pos; \n "
" posB.w = 0.f; \n "
" float4 c0local = convexShapes[shapeIndexA].m_localCenter; \n "
" float4 ornA = rigidBodies[bodyIndexA].m_quat; \n "
" float4 c0 = transform(&c0local, &posA, &ornA); \n "
" float4 c1local = convexShapes[shapeIndexB].m_localCenter; \n "
" float4 ornB =rigidBodies[bodyIndexB].m_quat; \n "
" float4 c1 = transform(&c1local,&posB,&ornB); \n "
" const float4 DeltaC2 = c0 - c1; \n "
" float4 sepNormal; \n "
" \n "
" bool sepA = findSeparatingAxis( &convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA, \n "
" posB,ornB, \n "
" DeltaC2, \n "
" vertices,uniqueEdges,faces, \n "
" indices,&sepNormal,&dmin); \n "
" hasSeparatingAxis[i] = 4; \n "
" if (!sepA) \n "
" { \n "
" hasSeparatingAxis[i] = 0; \n "
" } else \n "
" { \n "
" bool sepB = findSeparatingAxis( &convexShapes[shapeIndexB],&convexShapes[shapeIndexA],posB,ornB, \n "
" posA,ornA, \n "
" DeltaC2, \n "
" vertices,uniqueEdges,faces, \n "
" indices,&sepNormal,&dmin); \n "
" if (sepB) \n "
" { \n "
" dmins[i] = dmin; \n "
" hasSeparatingAxis[i] = 1; \n "
" separatingNormals[i] = sepNormal; \n "
" } \n "
" } \n "
" \n "
" } \n "
" } \n "
" __kernel void findSeparatingAxisEdgeEdgeKernel( __global const int4* pairs, \n "
" __global const BodyData* rigidBodies, \n "
" __global const btCollidableGpu* collidables, \n "
" __global const ConvexPolyhedronCL* convexShapes, \n "
" __global const float4* vertices, \n "
" __global const float4* uniqueEdges, \n "
" __global const btGpuFace* faces, \n "
" __global const int* indices, \n "
" __global btAabbCL* aabbs, \n "
" __global float4* separatingNormals, \n "
" __global int* hasSeparatingAxis, \n "
" __global float* dmins, \n "
" __global const float4* unitSphereDirections, \n "
" int numUnitSphereDirections, \n "
" int numPairs \n "
" ) \n "
" { \n "
" int i = get_global_id(0); \n "
" \n "
" if (i<numPairs) \n "
" { \n "
" if (hasSeparatingAxis[i]) \n "
" { \n "
" \n "
" int bodyIndexA = pairs[i].x; \n "
" int bodyIndexB = pairs[i].y; \n "
" \n "
" int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx; \n "
" int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx; \n "
" \n "
" int shapeIndexA = collidables[collidableIndexA].m_shapeIndex; \n "
" int shapeIndexB = collidables[collidableIndexB].m_shapeIndex; \n "
" \n "
" \n "
" int numFacesA = convexShapes[shapeIndexA].m_numFaces; \n "
" \n "
" float dmin = dmins[i]; \n "
" \n "
" float4 posA = rigidBodies[bodyIndexA].m_pos; \n "
" posA.w = 0.f; \n "
" float4 posB = rigidBodies[bodyIndexB].m_pos; \n "
" posB.w = 0.f; \n "
" float4 c0local = convexShapes[shapeIndexA].m_localCenter; \n "
" float4 ornA = rigidBodies[bodyIndexA].m_quat; \n "
" float4 c0 = transform(&c0local, &posA, &ornA); \n "
" float4 c1local = convexShapes[shapeIndexB].m_localCenter; \n "
" float4 ornB =rigidBodies[bodyIndexB].m_quat; \n "
" float4 c1 = transform(&c1local,&posB,&ornB); \n "
" const float4 DeltaC2 = c0 - c1; \n "
" float4 sepNormal = separatingNormals[i]; \n "
" \n "
" \n "
" \n "
" bool sepEE = false; \n "
" int numEdgeEdgeDirections = convexShapes[shapeIndexA].m_numUniqueEdges*convexShapes[shapeIndexB].m_numUniqueEdges; \n "
" if (numEdgeEdgeDirections<=numUnitSphereDirections) \n "
" { \n "
" sepEE = findSeparatingAxisEdgeEdge( &convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA, \n "
" posB,ornB, \n "
" DeltaC2, \n "
" vertices,uniqueEdges,faces, \n "
" indices,&sepNormal,&dmin); \n "
" \n "
" if (!sepEE) \n "
" { \n "
" hasSeparatingAxis[i] = 0; \n "
" } else \n "
" { \n "
" hasSeparatingAxis[i] = 1; \n "
" separatingNormals[i] = sepNormal; \n "
" } \n "
" } \n "
" /* \n "
" ///else case is a separate kernel, to make Mac OSX OpenCL compiler happy \n "
" else \n "
" { \n "
" sepEE = findSeparatingAxisUnitSphere(&convexShapes[shapeIndexA], &convexShapes[shapeIndexB],posA,ornA, \n "
" posB,ornB, \n "
" DeltaC2, \n "
" vertices,unitSphereDirections,numUnitSphereDirections, \n "
" &sepNormal,&dmin); \n "
" if (!sepEE) \n "
" { \n "
" hasSeparatingAxis[i] = 0; \n "
" } else \n "
" { \n "
" hasSeparatingAxis[i] = 1; \n "
" separatingNormals[i] = sepNormal; \n "
" } \n "
" } \n "
" */ \n "
" } //if (hasSeparatingAxis[i]) \n "
" }//(i<numPairs) \n "
" } \n "
" inline int findClippingFaces(const float4 separatingNormal, \n "
" const ConvexPolyhedronCL* hullA, \n "
" __global const ConvexPolyhedronCL* hullB, \n "
" const float4 posA, const Quaternion ornA,const float4 posB, const Quaternion ornB, \n "
" __global float4* worldVertsA1, \n "
" __global float4* worldNormalsA1, \n "
" __global float4* worldVertsB1, \n "
" int capacityWorldVerts, \n "
" const float minDist, float maxDist, \n "
" const float4* verticesA, \n "
" const btGpuFace* facesA, \n "
" const int* indicesA, \n "
" __global const float4* verticesB, \n "
" __global const btGpuFace* facesB, \n "
" __global const int* indicesB, \n "
" __global int4* clippingFaces, int pairIndex) \n "
" { \n "
" int numContactsOut = 0; \n "
" int numWorldVertsB1= 0; \n "
" \n "
" \n "
" int closestFaceB=0; \n "
" float dmax = -FLT_MAX; \n "
" \n "
" { \n "
" for(int face=0;face<hullB->m_numFaces;face++) \n "
" { \n "
" const float4 Normal = make_float4(facesB[hullB->m_faceOffset+face].m_plane.x, \n "
" facesB[hullB->m_faceOffset+face].m_plane.y, facesB[hullB->m_faceOffset+face].m_plane.z,0.f); \n "
" const float4 WorldNormal = qtRotate(ornB, Normal); \n "
" float d = dot3F4(WorldNormal,separatingNormal); \n "
" if (d > dmax) \n "
" { \n "
" dmax = d; \n "
" closestFaceB = face; \n "
" } \n "
" } \n "
" } \n "
" \n "
" { \n "
" const btGpuFace polyB = facesB[hullB->m_faceOffset+closestFaceB]; \n "
" int numVertices = polyB.m_numIndices; \n "
" if (numVertices>capacityWorldVerts) \n "
" numVertices = capacityWorldVerts; \n "
" \n "
" for(int e0=0;e0<numVertices;e0++) \n "
" { \n "
" if (e0<capacityWorldVerts) \n "
" { \n "
" const float4 b = verticesB[hullB->m_vertexOffset+indicesB[polyB.m_indexOffset+e0]]; \n "
" worldVertsB1[pairIndex*capacityWorldVerts+numWorldVertsB1++] = transform(&b,&posB,&ornB); \n "
" } \n "
" } \n "
" } \n "
" \n "
" int closestFaceA=0; \n "
" { \n "
" float dmin = FLT_MAX; \n "
" for(int face=0;face<hullA->m_numFaces;face++) \n "
" { \n "
" const float4 Normal = make_float4( \n "
" facesA[hullA->m_faceOffset+face].m_plane.x, \n "
" facesA[hullA->m_faceOffset+face].m_plane.y, \n "
" facesA[hullA->m_faceOffset+face].m_plane.z, \n "
" 0.f); \n "
" const float4 faceANormalWS = qtRotate(ornA,Normal); \n "
" \n "
" float d = dot3F4(faceANormalWS,separatingNormal); \n "
" if (d < dmin) \n "
" { \n "
" dmin = d; \n "
" closestFaceA = face; \n "
" worldNormalsA1[pairIndex] = faceANormalWS; \n "
" } \n "
" } \n "
" } \n "
" \n "
" int numVerticesA = facesA[hullA->m_faceOffset+closestFaceA].m_numIndices; \n "
" if (numVerticesA>capacityWorldVerts) \n "
" numVerticesA = capacityWorldVerts; \n "
" \n "
" for(int e0=0;e0<numVerticesA;e0++) \n "
" { \n "
" if (e0<capacityWorldVerts) \n "
" { \n "
" const float4 a = verticesA[hullA->m_vertexOffset+indicesA[facesA[hullA->m_faceOffset+closestFaceA].m_indexOffset+e0]]; \n "
" worldVertsA1[pairIndex*capacityWorldVerts+e0] = transform(&a, &posA,&ornA); \n "
" } \n "
" } \n "
" \n "
" clippingFaces[pairIndex].x = closestFaceA; \n "
" clippingFaces[pairIndex].y = closestFaceB; \n "
" clippingFaces[pairIndex].z = numVerticesA; \n "
" clippingFaces[pairIndex].w = numWorldVertsB1; \n "
" \n "
" \n "
" return numContactsOut; \n "
" } \n "
" // work-in-progress \n "
" __kernel void findConcaveSeparatingAxisKernel( __global int4* concavePairs, \n "
" __global const BodyData* rigidBodies, \n "
" __global const btCollidableGpu* collidables, \n "
" __global const ConvexPolyhedronCL* convexShapes, \n "
" __global const float4* vertices, \n "
" __global const float4* uniqueEdges, \n "
" __global const btGpuFace* faces, \n "
" __global const int* indices, \n "
" __global const btGpuChildShape* gpuChildShapes, \n "
" __global btAabbCL* aabbs, \n "
" __global float4* concaveSeparatingNormalsOut, \n "
" __global int* concaveHasSeparatingNormals, \n "
" __global int4* clippingFacesOut, \n "
" __global float4* worldVertsA1GPU, \n "
" __global float4* worldNormalsAGPU, \n "
" __global float4* worldVertsB1GPU, \n "
" int vertexFaceCapacity, \n "
" int numConcavePairs \n "
" ) \n "
" { \n "
" int i = get_global_id(0); \n "
" if (i>=numConcavePairs) \n "
" return; \n "
" concaveHasSeparatingNormals[i] = 0; \n "
" int pairIdx = i; \n "
" int bodyIndexA = concavePairs[i].x; \n "
" int bodyIndexB = concavePairs[i].y; \n "
" int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx; \n "
" int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx; \n "
" int shapeIndexA = collidables[collidableIndexA].m_shapeIndex; \n "
" int shapeIndexB = collidables[collidableIndexB].m_shapeIndex; \n "
" if (collidables[collidableIndexB].m_shapeType!=SHAPE_CONVEX_HULL&& \n "
" collidables[collidableIndexB].m_shapeType!=SHAPE_COMPOUND_OF_CONVEX_HULLS) \n "
" { \n "
" concavePairs[pairIdx].w = -1; \n "
" return; \n "
" } \n "
" int numFacesA = convexShapes[shapeIndexA].m_numFaces; \n "
" int numActualConcaveConvexTests = 0; \n "
" \n "
" int f = concavePairs[i].z; \n "
" \n "
" bool overlap = false; \n "
" \n "
" ConvexPolyhedronCL convexPolyhedronA; \n "
" //add 3 vertices of the triangle \n "
" convexPolyhedronA.m_numVertices = 3; \n "
" convexPolyhedronA.m_vertexOffset = 0; \n "
" float4 localCenter = make_float4(0.f,0.f,0.f,0.f); \n "
" btGpuFace face = faces[convexShapes[shapeIndexA].m_faceOffset+f]; \n "
" float4 triMinAabb, triMaxAabb; \n "
" btAabbCL triAabb; \n "
" triAabb.m_min = make_float4(1e30f,1e30f,1e30f,0.f); \n "
" triAabb.m_max = make_float4(-1e30f,-1e30f,-1e30f,0.f); \n "
" \n "
" float4 verticesA[3]; \n "
" for (int i=0;i<3;i++) \n "
" { \n "
" int index = indices[face.m_indexOffset+i]; \n "
" float4 vert = vertices[convexShapes[shapeIndexA].m_vertexOffset+index]; \n "
" verticesA[i] = vert; \n "
" localCenter += vert; \n "
" \n "
" triAabb.m_min = min(triAabb.m_min,vert); \n "
" triAabb.m_max = max(triAabb.m_max,vert); \n "
" } \n "
" overlap = true; \n "
" overlap = (triAabb.m_min.x > aabbs[bodyIndexB].m_max.x || triAabb.m_max.x < aabbs[bodyIndexB].m_min.x) ? false : overlap; \n "
" overlap = (triAabb.m_min.z > aabbs[bodyIndexB].m_max.z || triAabb.m_max.z < aabbs[bodyIndexB].m_min.z) ? false : overlap; \n "
" overlap = (triAabb.m_min.y > aabbs[bodyIndexB].m_max.y || triAabb.m_max.y < aabbs[bodyIndexB].m_min.y) ? false : overlap; \n "
" \n "
" if (overlap) \n "
" { \n "
" float dmin = FLT_MAX; \n "
" int hasSeparatingAxis=5; \n "
" float4 sepAxis=make_float4(1,2,3,4); \n "
" int localCC=0; \n "
" numActualConcaveConvexTests++; \n "
" //a triangle has 3 unique edges \n "
" convexPolyhedronA.m_numUniqueEdges = 3; \n "
" convexPolyhedronA.m_uniqueEdgesOffset = 0; \n "
" float4 uniqueEdgesA[3]; \n "
" \n "
" uniqueEdgesA[0] = (verticesA[1]-verticesA[0]); \n "
" uniqueEdgesA[1] = (verticesA[2]-verticesA[1]); \n "
" uniqueEdgesA[2] = (verticesA[0]-verticesA[2]); \n "
" convexPolyhedronA.m_faceOffset = 0; \n "
" \n "
" float4 normal = make_float4(face.m_plane.x,face.m_plane.y,face.m_plane.z,0.f); \n "
" \n "
" btGpuFace facesA[TRIANGLE_NUM_CONVEX_FACES]; \n "
" int indicesA[3+3+2+2+2]; \n "
" int curUsedIndices=0; \n "
" int fidx=0; \n "
" //front size of triangle \n "
" { \n "
" facesA[fidx].m_indexOffset=curUsedIndices; \n "
" indicesA[0] = 0; \n "
" indicesA[1] = 1; \n "
" indicesA[2] = 2; \n "
" curUsedIndices+=3; \n "
" float c = face.m_plane.w; \n "
" facesA[fidx].m_plane.x = normal.x; \n "
" facesA[fidx].m_plane.y = normal.y; \n "
" facesA[fidx].m_plane.z = normal.z; \n "
" facesA[fidx].m_plane.w = c; \n "
" facesA[fidx].m_numIndices=3; \n "
" } \n "
" fidx++; \n "
" //back size of triangle \n "
" { \n "
" facesA[fidx].m_indexOffset=curUsedIndices; \n "
" indicesA[3]=2; \n "
" indicesA[4]=1; \n "
" indicesA[5]=0; \n "
" curUsedIndices+=3; \n "
" float c = dot(normal,verticesA[0]); \n "
" float c1 = -face.m_plane.w; \n "
" facesA[fidx].m_plane.x = -normal.x; \n "
" facesA[fidx].m_plane.y = -normal.y; \n "
" facesA[fidx].m_plane.z = -normal.z; \n "
" facesA[fidx].m_plane.w = c; \n "
" facesA[fidx].m_numIndices=3; \n "
" } \n "
" fidx++; \n "
" bool addEdgePlanes = true; \n "
" if (addEdgePlanes) \n "
" { \n "
" int numVertices=3; \n "
" int prevVertex = numVertices-1; \n "
" for (int i=0;i<numVertices;i++) \n "
" { \n "
" float4 v0 = verticesA[i]; \n "
" float4 v1 = verticesA[prevVertex]; \n "
" \n "
" float4 edgeNormal = normalize(cross(normal,v1-v0)); \n "
" float c = -dot(edgeNormal,v0); \n "
" facesA[fidx].m_numIndices = 2; \n "
" facesA[fidx].m_indexOffset=curUsedIndices; \n "
" indicesA[curUsedIndices++]=i; \n "
" indicesA[curUsedIndices++]=prevVertex; \n "
" \n "
" facesA[fidx].m_plane.x = edgeNormal.x; \n "
" facesA[fidx].m_plane.y = edgeNormal.y; \n "
" facesA[fidx].m_plane.z = edgeNormal.z; \n "
" facesA[fidx].m_plane.w = c; \n "
" fidx++; \n "
" prevVertex = i; \n "
" } \n "
" } \n "
" convexPolyhedronA.m_numFaces = TRIANGLE_NUM_CONVEX_FACES; \n "
" convexPolyhedronA.m_localCenter = localCenter*(1.f/3.f); \n "
" float4 posA = rigidBodies[bodyIndexA].m_pos; \n "
" posA.w = 0.f; \n "
" float4 posB = rigidBodies[bodyIndexB].m_pos; \n "
" posB.w = 0.f; \n "
" float4 ornA = rigidBodies[bodyIndexA].m_quat; \n "
" float4 ornB =rigidBodies[bodyIndexB].m_quat; \n "
" \n "
" /////////////////// \n "
" ///compound shape support \n "
" if (collidables[collidableIndexB].m_shapeType==SHAPE_COMPOUND_OF_CONVEX_HULLS) \n "
" { \n "
" int compoundChild = concavePairs[pairIdx].w; \n "
" int childShapeIndexB = compoundChild;//collidables[collidableIndexB].m_shapeIndex+compoundChild; \n "
" int childColIndexB = gpuChildShapes[childShapeIndexB].m_shapeIndex; \n "
" float4 childPosB = gpuChildShapes[childShapeIndexB].m_childPosition; \n "
" float4 childOrnB = gpuChildShapes[childShapeIndexB].m_childOrientation; \n "
" float4 newPosB = transform(&childPosB,&posB,&ornB); \n "
" float4 newOrnB = qtMul(ornB,childOrnB); \n "
" posB = newPosB; \n "
" ornB = newOrnB; \n "
" shapeIndexB = collidables[childColIndexB].m_shapeIndex; \n "
" } \n "
" ////////////////// \n "
" float4 c0local = convexPolyhedronA.m_localCenter; \n "
" float4 c0 = transform(&c0local, &posA, &ornA); \n "
" float4 c1local = convexShapes[shapeIndexB].m_localCenter; \n "
" float4 c1 = transform(&c1local,&posB,&ornB); \n "
" const float4 DeltaC2 = c0 - c1; \n "
" bool sepA = findSeparatingAxisLocalA( &convexPolyhedronA, &convexShapes[shapeIndexB], \n "
" posA,ornA, \n "
" posB,ornB, \n "
" DeltaC2, \n "
" verticesA,uniqueEdgesA,facesA,indicesA, \n "
" vertices,uniqueEdges,faces,indices, \n "
" &sepAxis,&dmin); \n "
" hasSeparatingAxis = 4; \n "
" if (!sepA) \n "
" { \n "
" hasSeparatingAxis = 0; \n "
" } else \n "
" { \n "
" bool sepB = findSeparatingAxisLocalB( &convexShapes[shapeIndexB],&convexPolyhedronA, \n "
" posB,ornB, \n "
" posA,ornA, \n "
" DeltaC2, \n "
" vertices,uniqueEdges,faces,indices, \n "
" verticesA,uniqueEdgesA,facesA,indicesA, \n "
" &sepAxis,&dmin); \n "
" if (!sepB) \n "
" { \n "
" hasSeparatingAxis = 0; \n "
" } else \n "
" { \n "
" bool sepEE = findSeparatingAxisEdgeEdgeLocalA( &convexPolyhedronA, &convexShapes[shapeIndexB], \n "
" posA,ornA, \n "
" posB,ornB, \n "
" DeltaC2, \n "
" verticesA,uniqueEdgesA,facesA,indicesA, \n "
" vertices,uniqueEdges,faces,indices, \n "
" &sepAxis,&dmin); \n "
" \n "
" if (!sepEE) \n "
" { \n "
" hasSeparatingAxis = 0; \n "
" } else \n "
" { \n "
" hasSeparatingAxis = 1; \n "
" } \n "
" } \n "
" } \n "
" \n "
" if (hasSeparatingAxis) \n "
" { \n "
" sepAxis.w = dmin; \n "
" concaveSeparatingNormalsOut[pairIdx]=sepAxis; \n "
" concaveHasSeparatingNormals[i]=1; \n "
" float minDist = -1e30f; \n "
" float maxDist = 0.02f; \n "
" \n "
" findClippingFaces(sepAxis, \n "
" &convexPolyhedronA, \n "
" &convexShapes[shapeIndexB], \n "
" posA,ornA, \n "
" posB,ornB, \n "
" worldVertsA1GPU, \n "
" worldNormalsAGPU, \n "
" worldVertsB1GPU, \n "
" vertexFaceCapacity, \n "
" minDist, maxDist, \n "
" verticesA, \n "
" facesA, \n "
" indicesA, \n "
" vertices, \n "
" faces, \n "
" indices, \n "
" clippingFacesOut, pairIdx); \n "
" } else \n "
" { \n "
" //mark this pair as in-active \n "
" concavePairs[pairIdx].w = -1; \n "
" } \n "
" } \n "
" else \n "
" { \n "
" //mark this pair as in-active \n "
" concavePairs[pairIdx].w = -1; \n "
" } \n "
" \n "
" concavePairs[pairIdx].z = -1;//now z is used for existing/persistent contacts \n "
" } \n " ;
