src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/UpdateNormals.cl

   1 MSTRINGIFY(\r
   2 \r
   3 float length3(float4 a)\r
   4 {\r
   5         a.w = 0;\r
   6         return length(a);\r
   7 }\r
   8 \r
   9 float4 normalize3(float4 a)\r
  10 {\r
  11         a.w = 0;\r
  12         return normalize(a);\r
  13 }\r
  14 \r
  15 __kernel void \r
  16 ResetNormalsAndAreasKernel(\r
  17         const unsigned int numNodes,\r
  18         __global float4 * g_vertexNormals,\r
  19         __global float * g_vertexArea GUID_ARG)\r
  20 {\r
  21         if( get_global_id(0) < numNodes )\r
  22         {\r
  23                 g_vertexNormals[get_global_id(0)] = (float4)(0.0f, 0.0f, 0.0f, 0.0f);\r
  24                 g_vertexArea[get_global_id(0)]    = 0.0f;\r
  25         }\r
  26 }\r
  27 \r
  28 \r
  29 __kernel void \r
  30 UpdateSoftBodiesKernel(\r
  31         const unsigned int startFace,\r
  32         const unsigned int numFaces,\r
  33         __global int4 * g_triangleVertexIndexSet,\r
  34         __global float4 * g_vertexPositions,\r
  35         __global float4 * g_vertexNormals,\r
  36         __global float * g_vertexArea,\r
  37         __global float4 * g_triangleNormals,\r
  38         __global float * g_triangleArea GUID_ARG)\r
  39 {\r
  40         int faceID = get_global_id(0) + startFace;\r
  41         if( get_global_id(0) < numFaces )\r
  42         {               \r
  43                 int4 triangleIndexSet = g_triangleVertexIndexSet[ faceID ];\r
  44                 int nodeIndex0 = triangleIndexSet.x;\r
  45                 int nodeIndex1 = triangleIndexSet.y;\r
  46                 int nodeIndex2 = triangleIndexSet.z;\r
  47 \r
  48                 float4 node0 = g_vertexPositions[nodeIndex0];\r
  49                 float4 node1 = g_vertexPositions[nodeIndex1];\r
  50                 float4 node2 = g_vertexPositions[nodeIndex2];\r
  51                 float4 nodeNormal0 = g_vertexNormals[nodeIndex0];\r
  52                 float4 nodeNormal1 = g_vertexNormals[nodeIndex1];\r
  53                 float4 nodeNormal2 = g_vertexNormals[nodeIndex2];\r
  54                 float vertexArea0 = g_vertexArea[nodeIndex0];\r
  55                 float vertexArea1 = g_vertexArea[nodeIndex1];\r
  56                 float vertexArea2 = g_vertexArea[nodeIndex2];\r
  57                 \r
  58                 float4 vector0 = node1 - node0;\r
  59                 float4 vector1 = node2 - node0;\r
  60                 \r
  61                 float4 faceNormal = cross(vector0, vector1);\r
  62                 float triangleArea = length(faceNormal);\r
  63 \r
  64                 nodeNormal0 = nodeNormal0 + faceNormal;\r
  65                 nodeNormal1 = nodeNormal1 + faceNormal;\r
  66                 nodeNormal2 = nodeNormal2 + faceNormal;\r
  67                 vertexArea0 = vertexArea0 + triangleArea;\r
  68                 vertexArea1 = vertexArea1 + triangleArea;\r
  69                 vertexArea2 = vertexArea2 + triangleArea;\r
  70                 \r
  71                 g_triangleNormals[faceID] = normalize3(faceNormal);\r
  72                 g_vertexNormals[nodeIndex0] = nodeNormal0;\r
  73                 g_vertexNormals[nodeIndex1] = nodeNormal1;\r
  74                 g_vertexNormals[nodeIndex2] = nodeNormal2;\r
  75                 g_triangleArea[faceID] = triangleArea;\r
  76                 g_vertexArea[nodeIndex0] = vertexArea0;\r
  77                 g_vertexArea[nodeIndex1] = vertexArea1;\r
  78                 g_vertexArea[nodeIndex2] = vertexArea2;\r
  79         }\r
  80 }\r
  81 \r
  82 __kernel void \r
  83 NormalizeNormalsAndAreasKernel( \r
  84         const unsigned int numNodes,\r
  85         __global int * g_vertexTriangleCount,\r
  86         __global float4 * g_vertexNormals,\r
  87         __global float * g_vertexArea GUID_ARG)\r
  88 {\r
  89         if( get_global_id(0) < numNodes )\r
  90         {\r
  91                 float4 normal = g_vertexNormals[get_global_id(0)];\r
  92                 float area = g_vertexArea[get_global_id(0)];\r
  93                 int numTriangles = g_vertexTriangleCount[get_global_id(0)];\r
  94                 \r
  95                 float vectorLength = length3(normal);\r
  96                 \r
  97                 g_vertexNormals[get_global_id(0)] = normalize3(normal);\r
  98                 g_vertexArea[get_global_id(0)] = area/(float)(numTriangles);\r
  99         }\r
 100 }\r
 101 \r
 102 );\r