← Back to branch summary

~ubuntu-branches/ubuntu/vivid/emscripten/vivid

« back to all changes in this revision

Viewing changes to tests/bullet/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/UpdateNormals.cl

  • Committer: Package Import Robot
  • Author(s): Sylvestre Ledru
  • Date: 2013-05-02 13:11:51 UTC
  • Revision ID: package-import@ubuntu.com-20130502131151-q8dvteqr1ef2x7xz
Tags: upstream-1.4.1~20130504~adb56cb
Import upstream version 1.4.1~20130504~adb56cb

Show diffs side-by-side

added added

removed removed

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