aboutsummaryrefslogtreecommitdiff
path: root/tests/bullet/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC10/SolvePositions.cl
blob: fe7aec66eebcc8be2ecce82eb06c2dd780348adf (plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
MSTRINGIFY(


float mydot3(float4 a, float4 b)
{
   return a.x*b.x + a.y*b.y + a.z*b.z;
}


__kernel void 
SolvePositionsFromLinksKernel( 
	const int startLink,
	const int numLinks,
	const float kst,
	const float ti,
	__global int2 * g_linksVertexIndices,
	__global float * g_linksMassLSC,
	__global float * g_linksRestLengthSquared,
	__global float * g_verticesInverseMass,
	__global float4 * g_vertexPositions GUID_ARG)
	
{
	int linkID = get_global_id(0) + startLink;
	if( get_global_id(0) < numLinks )
	{	
		float massLSC = g_linksMassLSC[linkID];
		float restLengthSquared = g_linksRestLengthSquared[linkID];
		
		if( massLSC > 0.0f )
		{		
			int2 nodeIndices = g_linksVertexIndices[linkID];
			int node0 = nodeIndices.x;
			int node1 = nodeIndices.y;
			
			float4 position0 = g_vertexPositions[node0];
			float4 position1 = g_vertexPositions[node1];

			float inverseMass0 = g_verticesInverseMass[node0];
			float inverseMass1 = g_verticesInverseMass[node1]; 

			float4 del = position1 - position0;
			float len  = mydot3(del, del);
			float k    = ((restLengthSquared - len)/(massLSC*(restLengthSquared+len)))*kst;
			position0 = position0 - del*(k*inverseMass0);
			position1 = position1 + del*(k*inverseMass1);

			g_vertexPositions[node0] = position0;
			g_vertexPositions[node1] = position1;

		}
	}
}

);