summaryrefslogtreecommitdiff
path: root/tests/bullet/src/BulletMultiThreaded/GpuSoftBodySolvers/OpenCL/OpenCLC/Integrate.cl
blob: 4a2c9f2fcc2faa3463b5c51d909018a85020f1d1 (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
MSTRINGIFY(

// Node indices for each link

//#define float3 float4

__kernel void
IntegrateKernel( 
	const int numNodes,
	const float solverdt,
	__global float * g_vertexInverseMasses,
	__global float4 * g_vertexPositions,
	__global float4 * g_vertexVelocity,
	__global float4 * g_vertexPreviousPositions,
	__global float4 * g_vertexForceAccumulator)
{
	int nodeID = get_global_id(0);
	if( nodeID < numNodes )
	{	
		float3 position   = g_vertexPositions[nodeID].xyz;
		float3 velocity   = g_vertexVelocity[nodeID].xyz;
		float3 force      = g_vertexForceAccumulator[nodeID].xyz;
		float inverseMass = g_vertexInverseMasses[nodeID];
		
		g_vertexPreviousPositions[nodeID] = (float4)(position, 0.f);
		velocity += force * inverseMass * solverdt;
		position += velocity * solverdt;
		
		g_vertexForceAccumulator[nodeID] = (float4)(0.f, 0.f, 0.f, 0.0f);
		g_vertexPositions[nodeID]        = (float4)(position, 0.f);
		g_vertexVelocity[nodeID]         = (float4)(velocity, 0.f);	
	}
}

);