dali-physics/third-party/bullet3/src/Bullet3OpenCL/Raycast/kernels/rayCastKernels.h

   1 //this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
   2 static const char* rayCastKernelCL =
   3         "#define SHAPE_CONVEX_HULL 3\n"
   4         "#define SHAPE_PLANE 4\n"
   5         "#define SHAPE_CONCAVE_TRIMESH 5\n"
   6         "#define SHAPE_COMPOUND_OF_CONVEX_HULLS 6\n"
   7         "#define SHAPE_SPHERE 7\n"
   8         "typedef struct\n"
   9         "{\n"
  10         "       float4 m_from;\n"
  11         "       float4 m_to;\n"
  12         "} b3RayInfo;\n"
  13         "typedef struct\n"
  14         "{\n"
  15         "       float m_hitFraction;\n"
  16         "       int     m_hitResult0;\n"
  17         "       int     m_hitResult1;\n"
  18         "       int     m_hitResult2;\n"
  19         "       float4  m_hitPoint;\n"
  20         "       float4  m_hitNormal;\n"
  21         "} b3RayHit;\n"
  22         "typedef struct\n"
  23         "{\n"
  24         "       float4 m_pos;\n"
  25         "       float4 m_quat;\n"
  26         "       float4 m_linVel;\n"
  27         "       float4 m_angVel;\n"
  28         "       unsigned int m_collidableIdx;\n"
  29         "       float m_invMass;\n"
  30         "       float m_restituitionCoeff;\n"
  31         "       float m_frictionCoeff;\n"
  32         "} Body;\n"
  33         "typedef struct Collidable\n"
  34         "{\n"
  35         "       union {\n"
  36         "               int m_numChildShapes;\n"
  37         "               int m_bvhIndex;\n"
  38         "       };\n"
  39         "       float m_radius;\n"
  40         "       int m_shapeType;\n"
  41         "       int m_shapeIndex;\n"
  42         "} Collidable;\n"
  43         "typedef struct  \n"
  44         "{\n"
  45         "       float4          m_localCenter;\n"
  46         "       float4          m_extents;\n"
  47         "       float4          mC;\n"
  48         "       float4          mE;\n"
  49         "       float                   m_radius;\n"
  50         "       int     m_faceOffset;\n"
  51         "       int m_numFaces;\n"
  52         "       int     m_numVertices;\n"
  53         "       int m_vertexOffset;\n"
  54         "       int     m_uniqueEdgesOffset;\n"
  55         "       int     m_numUniqueEdges;\n"
  56         "       int m_unused;\n"
  57         "} ConvexPolyhedronCL;\n"
  58         "typedef struct\n"
  59         "{\n"
  60         "       float4 m_plane;\n"
  61         "       int m_indexOffset;\n"
  62         "       int m_numIndices;\n"
  63         "} b3GpuFace;\n"
  64         "///////////////////////////////////////\n"
  65         "//     Quaternion\n"
  66         "///////////////////////////////////////\n"
  67         "typedef float4 Quaternion;\n"
  68         "__inline\n"
  69         "       Quaternion qtMul(Quaternion a, Quaternion b);\n"
  70         "__inline\n"
  71         "       Quaternion qtNormalize(Quaternion in);\n"
  72         "__inline\n"
  73         "       Quaternion qtInvert(Quaternion q);\n"
  74         "__inline\n"
  75         "       float dot3F4(float4 a, float4 b)\n"
  76         "{\n"
  77         "       float4 a1 = (float4)(a.xyz,0.f);\n"
  78         "       float4 b1 = (float4)(b.xyz,0.f);\n"
  79         "       return dot(a1, b1);\n"
  80         "}\n"
  81         "__inline\n"
  82         "       Quaternion qtMul(Quaternion a, Quaternion b)\n"
  83         "{\n"
  84         "       Quaternion ans;\n"
  85         "       ans = cross( a, b );\n"
  86         "       ans += a.w*b+b.w*a;\n"
  87         "       //      ans.w = a.w*b.w - (a.x*b.x+a.y*b.y+a.z*b.z);\n"
  88         "       ans.w = a.w*b.w - dot3F4(a, b);\n"
  89         "       return ans;\n"
  90         "}\n"
  91         "__inline\n"
  92         "       Quaternion qtNormalize(Quaternion in)\n"
  93         "{\n"
  94         "       return fast_normalize(in);\n"
  95         "       //      in /= length( in );\n"
  96         "       //      return in;\n"
  97         "}\n"
  98         "__inline\n"
  99         "       float4 qtRotate(Quaternion q, float4 vec)\n"
 100         "{\n"
 101         "       Quaternion qInv = qtInvert( q );\n"
 102         "       float4 vcpy = vec;\n"
 103         "       vcpy.w = 0.f;\n"
 104         "       float4 out = qtMul(q,vcpy);\n"
 105         "       out = qtMul(out,qInv);\n"
 106         "       return out;\n"
 107         "}\n"
 108         "__inline\n"
 109         "       Quaternion qtInvert(Quaternion q)\n"
 110         "{\n"
 111         "       return (Quaternion)(-q.xyz, q.w);\n"
 112         "}\n"
 113         "__inline\n"
 114         "       float4 qtInvRotate(const Quaternion q, float4 vec)\n"
 115         "{\n"
 116         "       return qtRotate( qtInvert( q ), vec );\n"
 117         "}\n"
 118         "void   trInverse(float4 translationIn, Quaternion orientationIn,\n"
 119         "       float4* translationOut, Quaternion* orientationOut)\n"
 120         "{\n"
 121         "       *orientationOut = qtInvert(orientationIn);\n"
 122         "       *translationOut = qtRotate(*orientationOut, -translationIn);\n"
 123         "}\n"
 124         "bool rayConvex(float4 rayFromLocal, float4 rayToLocal, int numFaces, int faceOffset,\n"
 125         "       __global const b3GpuFace* faces, float* hitFraction, float4* hitNormal)\n"
 126         "{\n"
 127         "       rayFromLocal.w = 0.f;\n"
 128         "       rayToLocal.w = 0.f;\n"
 129         "       bool result = true;\n"
 130         "       float exitFraction = hitFraction[0];\n"
 131         "       float enterFraction = -0.3f;\n"
 132         "       float4 curHitNormal = (float4)(0,0,0,0);\n"
 133         "       for (int i=0;i<numFaces && result;i++)\n"
 134         "       {\n"
 135         "               b3GpuFace face = faces[faceOffset+i];\n"
 136         "               float fromPlaneDist = dot(rayFromLocal,face.m_plane)+face.m_plane.w;\n"
 137         "               float toPlaneDist = dot(rayToLocal,face.m_plane)+face.m_plane.w;\n"
 138         "               if (fromPlaneDist<0.f)\n"
 139         "               {\n"
 140         "                       if (toPlaneDist >= 0.f)\n"
 141         "                       {\n"
 142         "                               float fraction = fromPlaneDist / (fromPlaneDist-toPlaneDist);\n"
 143         "                               if (exitFraction>fraction)\n"
 144         "                               {\n"
 145         "                                       exitFraction = fraction;\n"
 146         "                               }\n"
 147         "                       }                       \n"
 148         "               } else\n"
 149         "               {\n"
 150         "                       if (toPlaneDist<0.f)\n"
 151         "                       {\n"
 152         "                               float fraction = fromPlaneDist / (fromPlaneDist-toPlaneDist);\n"
 153         "                               if (enterFraction <= fraction)\n"
 154         "                               {\n"
 155         "                                       enterFraction = fraction;\n"
 156         "                                       curHitNormal = face.m_plane;\n"
 157         "                                       curHitNormal.w = 0.f;\n"
 158         "                               }\n"
 159         "                       } else\n"
 160         "                       {\n"
 161         "                               result = false;\n"
 162         "                       }\n"
 163         "               }\n"
 164         "               if (exitFraction <= enterFraction)\n"
 165         "                       result = false;\n"
 166         "       }\n"
 167         "       if (enterFraction < 0.f)\n"
 168         "       {\n"
 169         "               result = false;\n"
 170         "       }\n"
 171         "       if (result)\n"
 172         "       {       \n"
 173         "               hitFraction[0] = enterFraction;\n"
 174         "               hitNormal[0] = curHitNormal;\n"
 175         "       }\n"
 176         "       return result;\n"
 177         "}\n"
 178         "bool sphere_intersect(float4 spherePos,  float radius, float4 rayFrom, float4 rayTo, float* hitFraction)\n"
 179         "{\n"
 180         "       float4 rs = rayFrom - spherePos;\n"
 181         "       rs.w = 0.f;\n"
 182         "       float4 rayDir = rayTo-rayFrom;\n"
 183         "       rayDir.w = 0.f;\n"
 184         "       float A = dot(rayDir,rayDir);\n"
 185         "       float B = dot(rs, rayDir);\n"
 186         "       float C = dot(rs, rs) - (radius * radius);\n"
 187         "       float D = B * B - A*C;\n"
 188         "       if (D > 0.0f)\n"
 189         "       {\n"
 190         "               float t = (-B - sqrt(D))/A;\n"
 191         "               if ( (t >= 0.0f) && (t < (*hitFraction)) )\n"
 192         "               {\n"
 193         "                       *hitFraction = t;\n"
 194         "                       return true;\n"
 195         "               }\n"
 196         "       }\n"
 197         "       return false;\n"
 198         "}\n"
 199         "float4 setInterpolate3(float4 from, float4 to, float t)\n"
 200         "{\n"
 201         "       float s = 1.0f - t;\n"
 202         "       float4 result;\n"
 203         "       result = s * from + t * to;\n"
 204         "       result.w = 0.f; \n"
 205         "       return result;  \n"
 206         "}\n"
 207         "__kernel void rayCastKernel(  \n"
 208         "       int numRays, \n"
 209         "       const __global b3RayInfo* rays, \n"
 210         "       __global b3RayHit* hitResults, \n"
 211         "       const int numBodies, \n"
 212         "       __global Body* bodies,\n"
 213         "       __global Collidable* collidables,\n"
 214         "       __global const b3GpuFace* faces,\n"
 215         "       __global const ConvexPolyhedronCL* convexShapes )\n"
 216         "{\n"
 217         "       int i = get_global_id(0);\n"
 218         "       if (i>=numRays)\n"
 219         "               return;\n"
 220         "       hitResults[i].m_hitFraction = 1.f;\n"
 221         "       float4 rayFrom = rays[i].m_from;\n"
 222         "       float4 rayTo = rays[i].m_to;\n"
 223         "       float hitFraction = 1.f;\n"
 224         "       float4 hitPoint;\n"
 225         "       float4 hitNormal;\n"
 226         "       int hitBodyIndex= -1;\n"
 227         "       int cachedCollidableIndex = -1;\n"
 228         "       Collidable cachedCollidable;\n"
 229         "       for (int b=0;b<numBodies;b++)\n"
 230         "       {\n"
 231         "               if (hitResults[i].m_hitResult2==b)\n"
 232         "                       continue;\n"
 233         "               Body body = bodies[b];\n"
 234         "               float4 pos = body.m_pos;\n"
 235         "               float4 orn = body.m_quat;\n"
 236         "               if (cachedCollidableIndex != body.m_collidableIdx)\n"
 237         "               {\n"
 238         "                       cachedCollidableIndex = body.m_collidableIdx;\n"
 239         "                       cachedCollidable = collidables[cachedCollidableIndex];\n"
 240         "               }\n"
 241         "               if (cachedCollidable.m_shapeType == SHAPE_CONVEX_HULL)\n"
 242         "               {\n"
 243         "                       float4 invPos = (float4)(0,0,0,0);\n"
 244         "                       float4 invOrn = (float4)(0,0,0,0);\n"
 245         "                       float4 rayFromLocal = (float4)(0,0,0,0);\n"
 246         "                       float4 rayToLocal = (float4)(0,0,0,0);\n"
 247         "                       invOrn = qtInvert(orn);\n"
 248         "                       invPos = qtRotate(invOrn, -pos);\n"
 249         "                       rayFromLocal = qtRotate( invOrn, rayFrom ) + invPos;\n"
 250         "                       rayToLocal = qtRotate( invOrn, rayTo) + invPos;\n"
 251         "                       rayFromLocal.w = 0.f;\n"
 252         "                       rayToLocal.w = 0.f;\n"
 253         "                       int numFaces = convexShapes[cachedCollidable.m_shapeIndex].m_numFaces;\n"
 254         "                       int faceOffset = convexShapes[cachedCollidable.m_shapeIndex].m_faceOffset;\n"
 255         "                       if (numFaces)\n"
 256         "                       {\n"
 257         "                               if (rayConvex(rayFromLocal, rayToLocal, numFaces, faceOffset,faces, &hitFraction, &hitNormal))\n"
 258         "                               {\n"
 259         "                                       hitBodyIndex = b;\n"
 260         "                                       \n"
 261         "                               }\n"
 262         "                       }\n"
 263         "               }\n"
 264         "               if (cachedCollidable.m_shapeType == SHAPE_SPHERE)\n"
 265         "               {\n"
 266         "                       float radius = cachedCollidable.m_radius;\n"
 267         "               \n"
 268         "                       if (sphere_intersect(pos,  radius, rayFrom, rayTo, &hitFraction))\n"
 269         "                       {\n"
 270         "                               hitBodyIndex = b;\n"
 271         "                               hitNormal = (float4) (hitPoint-bodies[b].m_pos);\n"
 272         "                       }\n"
 273         "               }\n"
 274         "       }\n"
 275         "       if (hitBodyIndex>=0)\n"
 276         "       {\n"
 277         "               hitPoint = setInterpolate3(rayFrom, rayTo,hitFraction);\n"
 278         "               hitResults[i].m_hitFraction = hitFraction;\n"
 279         "               hitResults[i].m_hitPoint = hitPoint;\n"
 280         "               hitResults[i].m_hitNormal = normalize(hitNormal);\n"
 281         "               hitResults[i].m_hitResult0 = hitBodyIndex;\n"
 282         "       }\n"
 283         "}\n"
 284         "__kernel void findRayRigidPairIndexRanges(__global int2* rayRigidPairs, \n"
 285         "                                                                                       __global int* out_firstRayRigidPairIndexPerRay,\n"
 286         "                                                                                       __global int* out_numRayRigidPairsPerRay,\n"
 287         "                                                                                       int numRayRigidPairs)\n"
 288         "{\n"
 289         "       int rayRigidPairIndex = get_global_id(0);\n"
 290         "       if (rayRigidPairIndex >= numRayRigidPairs) return;\n"
 291         "       \n"
 292         "       int rayIndex = rayRigidPairs[rayRigidPairIndex].x;\n"
 293         "       \n"
 294         "       atomic_min(&out_firstRayRigidPairIndexPerRay[rayIndex], rayRigidPairIndex);\n"
 295         "       atomic_inc(&out_numRayRigidPairsPerRay[rayIndex]);\n"
 296         "}\n"
 297         "__kernel void rayCastPairsKernel(const __global b3RayInfo* rays, \n"
 298         "                                                               __global b3RayHit* hitResults, \n"
 299         "                                                               __global int* firstRayRigidPairIndexPerRay,\n"
 300         "                                                               __global int* numRayRigidPairsPerRay,\n"
 301         "                                                                       \n"
 302         "                                                               __global Body* bodies,\n"
 303         "                                                               __global Collidable* collidables,\n"
 304         "                                                               __global const b3GpuFace* faces,\n"
 305         "                                                               __global const ConvexPolyhedronCL* convexShapes,\n"
 306         "                                                               \n"
 307         "                                                               __global int2* rayRigidPairs,\n"
 308         "                                                               int numRays)\n"
 309         "{\n"
 310         "       int i = get_global_id(0);\n"
 311         "       if (i >= numRays) return;\n"
 312         "       \n"
 313         "       float4 rayFrom = rays[i].m_from;\n"
 314         "       float4 rayTo = rays[i].m_to;\n"
 315         "               \n"
 316         "       hitResults[i].m_hitFraction = 1.f;\n"
 317         "               \n"
 318         "       float hitFraction = 1.f;\n"
 319         "       float4 hitPoint;\n"
 320         "       float4 hitNormal;\n"
 321         "       int hitBodyIndex = -1;\n"
 322         "               \n"
 323         "       //\n"
 324         "       for(int pair = 0; pair < numRayRigidPairsPerRay[i]; ++pair)\n"
 325         "       {\n"
 326         "               int rayRigidPairIndex = pair + firstRayRigidPairIndexPerRay[i];\n"
 327         "               int b = rayRigidPairs[rayRigidPairIndex].y;\n"
 328         "               \n"
 329         "               if (hitResults[i].m_hitResult2 == b) continue;\n"
 330         "               \n"
 331         "               Body body = bodies[b];\n"
 332         "               Collidable rigidCollidable = collidables[body.m_collidableIdx];\n"
 333         "               \n"
 334         "               float4 pos = body.m_pos;\n"
 335         "               float4 orn = body.m_quat;\n"
 336         "               \n"
 337         "               if (rigidCollidable.m_shapeType == SHAPE_CONVEX_HULL)\n"
 338         "               {\n"
 339         "                       float4 invPos = (float4)(0,0,0,0);\n"
 340         "                       float4 invOrn = (float4)(0,0,0,0);\n"
 341         "                       float4 rayFromLocal = (float4)(0,0,0,0);\n"
 342         "                       float4 rayToLocal = (float4)(0,0,0,0);\n"
 343         "                       invOrn = qtInvert(orn);\n"
 344         "                       invPos = qtRotate(invOrn, -pos);\n"
 345         "                       rayFromLocal = qtRotate( invOrn, rayFrom ) + invPos;\n"
 346         "                       rayToLocal = qtRotate( invOrn, rayTo) + invPos;\n"
 347         "                       rayFromLocal.w = 0.f;\n"
 348         "                       rayToLocal.w = 0.f;\n"
 349         "                       int numFaces = convexShapes[rigidCollidable.m_shapeIndex].m_numFaces;\n"
 350         "                       int faceOffset = convexShapes[rigidCollidable.m_shapeIndex].m_faceOffset;\n"
 351         "                       \n"
 352         "                       if (numFaces && rayConvex(rayFromLocal, rayToLocal, numFaces, faceOffset,faces, &hitFraction, &hitNormal))\n"
 353         "                       {\n"
 354         "                               hitBodyIndex = b;\n"
 355         "                               hitPoint = setInterpolate3(rayFrom, rayTo, hitFraction);\n"
 356         "                       }\n"
 357         "               }\n"
 358         "               \n"
 359         "               if (rigidCollidable.m_shapeType == SHAPE_SPHERE)\n"
 360         "               {\n"
 361         "                       float radius = rigidCollidable.m_radius;\n"
 362         "               \n"
 363         "                       if (sphere_intersect(pos, radius, rayFrom, rayTo, &hitFraction))\n"
 364         "                       {\n"
 365         "                               hitBodyIndex = b;\n"
 366         "                               hitPoint = setInterpolate3(rayFrom, rayTo, hitFraction);\n"
 367         "                               hitNormal = (float4) (hitPoint - bodies[b].m_pos);\n"
 368         "                       }\n"
 369         "               }\n"
 370         "       }\n"
 371         "       \n"
 372         "       if (hitBodyIndex >= 0)\n"
 373         "       {\n"
 374         "               hitResults[i].m_hitFraction = hitFraction;\n"
 375         "               hitResults[i].m_hitPoint = hitPoint;\n"
 376         "               hitResults[i].m_hitNormal = normalize(hitNormal);\n"
 377         "               hitResults[i].m_hitResult0 = hitBodyIndex;\n"
 378         "       }\n"
 379         "       \n"
 380         "}\n";