← 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/SolveCollisionsAndUpdateVelocitiesSIMDBatched.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/SolveCollisionsAndUpdateVelocitiesSIMDBatched.cl
 
1
MSTRINGIFY(
 
2
float mydot3(float4 a, float4 b)
 
3
{
 
4
   return a.x*b.x + a.y*b.y + a.z*b.z;
 
5
}
 
6
 
 
7
 
 
8
typedef struct 
 
9
{
 
10
        int firstObject;
 
11
        int endObject;
 
12
} CollisionObjectIndices;
 
13
 
 
14
typedef struct 
 
15
{
 
16
        float4 shapeTransform[4]; // column major 4x4 matrix
 
17
        float4 linearVelocity;
 
18
        float4 angularVelocity;
 
19
 
 
20
        int softBodyIdentifier;
 
21
        int collisionShapeType;
 
22
        
 
23
 
 
24
        // Shape information
 
25
        // Compressed from the union
 
26
        float radius;
 
27
        float halfHeight;
 
28
        int upAxis;
 
29
                
 
30
        float margin;
 
31
        float friction;
 
32
 
 
33
        int padding0;
 
34
        
 
35
} CollisionShapeDescription;
 
36
 
 
37
// From btBroadphaseProxy.h
 
38
__constant int CAPSULE_SHAPE_PROXYTYPE = 10;
 
39
 
 
40
 
 
41
/* Multiply column-major matrix against vector */
 
42
float4 matrixVectorMul( float4 matrix[4], float4 vector )
 
43
{
 
44
        float4 returnVector;
 
45
        float4 row0 = (float4)(matrix[0].x, matrix[1].x, matrix[2].x, matrix[3].x);
 
46
        float4 row1 = (float4)(matrix[0].y, matrix[1].y, matrix[2].y, matrix[3].y);
 
47
        float4 row2 = (float4)(matrix[0].z, matrix[1].z, matrix[2].z, matrix[3].z);
 
48
        float4 row3 = (float4)(matrix[0].w, matrix[1].w, matrix[2].w, matrix[3].w);
 
49
        returnVector.x = dot(row0, vector);
 
50
        returnVector.y = dot(row1, vector);
 
51
        returnVector.z = dot(row2, vector);
 
52
        returnVector.w = dot(row3, vector);
 
53
        return returnVector;
 
54
}
 
55
 
 
56
__kernel void 
 
57
SolveCollisionsAndUpdateVelocitiesKernel( 
 
58
        const int numNodes,
 
59
        const float isolverdt,
 
60
        __global int *g_vertexClothIdentifier,
 
61
        __global float4 *g_vertexPreviousPositions,
 
62
        __global float * g_perClothFriction,
 
63
        __global float * g_clothDampingFactor,
 
64
        __global CollisionObjectIndices * g_perClothCollisionObjectIndices,
 
65
        __global CollisionShapeDescription * g_collisionObjectDetails,
 
66
        __global float4 * g_vertexForces,
 
67
        __global float4 *g_vertexVelocities,
 
68
        __global float4 *g_vertexPositions,
 
69
        __local CollisionShapeDescription *localCollisionShapes)
 
70
{
 
71
        int nodeID = get_global_id(0);
 
72
        float4 forceOnVertex = (float4)(0.f, 0.f, 0.f, 0.f);
 
73
 
 
74
        int clothIdentifier = g_vertexClothIdentifier[nodeID];
 
75
 
 
76
        // Abort if this is not a valid cloth
 
77
        if( clothIdentifier < 0 )
 
78
                return;
 
79
        
 
80
 
 
81
        float4 position = (float4)(g_vertexPositions[nodeID].xyz, 1.f);
 
82
        float4 previousPosition = (float4)(g_vertexPreviousPositions[nodeID].xyz, 1.f);
 
83
        float clothFriction = g_perClothFriction[clothIdentifier];
 
84
        float dampingFactor = g_clothDampingFactor[clothIdentifier];
 
85
        float velocityCoefficient = (1.f - dampingFactor);              
 
86
        
 
87
        // Update velocity      
 
88
        float4 difference = position - previousPosition;
 
89
        float4 velocity = difference*velocityCoefficient*isolverdt;                     
 
90
        CollisionObjectIndices collisionObjectIndices = g_perClothCollisionObjectIndices[clothIdentifier];
 
91
        
 
92
        int numObjects = collisionObjectIndices.endObject - collisionObjectIndices.firstObject;
 
93
        if( numObjects > 0 )
 
94
        {
 
95
                // We have some possible collisions to deal with
 
96
                
 
97
                // First load all of the collision objects into LDS
 
98
                int numObjects = collisionObjectIndices.endObject - collisionObjectIndices.firstObject;
 
99
                if( get_local_id(0) < numObjects )
 
100
                {
 
101
                        localCollisionShapes[get_local_id(0)] = g_collisionObjectDetails[ collisionObjectIndices.firstObject + get_local_id(0) ];
 
102
                }
 
103
        }
 
104
 
 
105
        // Safe as the vertices are padded so that not more than one soft body is in a group
 
106
        barrier(CLK_LOCAL_MEM_FENCE);
 
107
 
 
108
        // Annoyingly, even though I know the flow control is not varying, the compiler will not let me skip this
 
109
        if( numObjects > 0 )
 
110
        {
 
111
                
 
112
                
 
113
                // We have some possible collisions to deal with
 
114
                for( int collision = 0; collision < numObjects; ++collision )
 
115
                {
 
116
                        //CollisionShapeDescription shapeDescription = localCollisionShapes[collision];
 
117
                        float colliderFriction = localCollisionShapes[collision].friction;
 
118
                
 
119
                        if( localCollisionShapes[collision].collisionShapeType == CAPSULE_SHAPE_PROXYTYPE )
 
120
                        {
 
121
                                // Colliding with a capsule
 
122
 
 
123
                                float capsuleHalfHeight = localCollisionShapes[collision].halfHeight;
 
124
                                float capsuleRadius = localCollisionShapes[collision].radius;
 
125
                                float capsuleMargin = localCollisionShapes[collision].margin;
 
126
                                int capsuleupAxis = localCollisionShapes[collision].upAxis;
 
127
 
 
128
                                float4 worldTransform[4];
 
129
                                worldTransform[0] = localCollisionShapes[collision].shapeTransform[0];
 
130
                                worldTransform[1] = localCollisionShapes[collision].shapeTransform[1];
 
131
                                worldTransform[2] = localCollisionShapes[collision].shapeTransform[2];
 
132
                                worldTransform[3] = localCollisionShapes[collision].shapeTransform[3];
 
133
 
 
134
                                //float4 c1 = (float4)(0.f, -capsuleHalfHeight, 0.f, 1.f); 
 
135
                                //float4 c2 = (float4)(0.f, +capsuleHalfHeight, 0.f, 1.f);
 
136
                                // Correctly define capsule centerline vector 
 
137
                                float4 c1 = (float4)(0.f, 0.f, 0.f, 1.f); 
 
138
                                float4 c2 = (float4)(0.f, 0.f, 0.f, 1.f);
 
139
                                c1.x = select( 0.f, -capsuleHalfHeight, capsuleupAxis == 0 );
 
140
                                c1.y = select( 0.f, -capsuleHalfHeight, capsuleupAxis == 1 );
 
141
                                c1.z = select( 0.f, -capsuleHalfHeight, capsuleupAxis == 2 );
 
142
                                c2.x = -c1.x;
 
143
                                c2.y = -c1.y;
 
144
                                c2.z = -c1.z;
 
145
 
 
146
                                float4 worldC1 = matrixVectorMul(worldTransform, c1);
 
147
                                float4 worldC2 = matrixVectorMul(worldTransform, c2);
 
148
                                float4 segment = (worldC2 - worldC1);
 
149
 
 
150
 
 
151
                                // compute distance of tangent to vertex along line segment in capsule
 
152
                                float distanceAlongSegment = -( mydot3( (worldC1 - position), segment ) / mydot3(segment, segment) );
 
153
 
 
154
                                float4 closestPoint = (worldC1 + (float4)(segment * distanceAlongSegment));
 
155
                                float distanceFromLine = length(position - closestPoint);
 
156
                                float distanceFromC1 = length(worldC1 - position);
 
157
                                float distanceFromC2 = length(worldC2 - position);
 
158
                                        
 
159
                                // Final distance from collision, point to push from, direction to push in
 
160
                                // for impulse force
 
161
                                float dist;
 
162
                                float4 normalVector;
 
163
                                if( distanceAlongSegment < 0 )
 
164
                                {
 
165
                                        dist = distanceFromC1;
 
166
                                        normalVector = normalize(position - worldC1);
 
167
                                } else if( distanceAlongSegment > 1.f ) {
 
168
                                        dist = distanceFromC2;
 
169
                                        normalVector = normalize(position - worldC2);   
 
170
                                } else {
 
171
                                        dist = distanceFromLine;
 
172
                                        normalVector = normalize(position - closestPoint);
 
173
                                }
 
174
                                                
 
175
                                float4 colliderLinearVelocity = localCollisionShapes[collision].linearVelocity;
 
176
                                float4 colliderAngularVelocity = localCollisionShapes[collision].angularVelocity;
 
177
                                float4 velocityOfSurfacePoint = colliderLinearVelocity + cross(colliderAngularVelocity, position - (float4)(worldTransform[0].w, worldTransform[1].w, worldTransform[2].w, 0.f));
 
178
 
 
179
                                float minDistance = capsuleRadius + capsuleMargin;
 
180
                                        
 
181
                                // In case of no collision, this is the value of velocity
 
182
                                velocity = (position - previousPosition) * velocityCoefficient * isolverdt;
 
183
                                        
 
184
                                        
 
185
                                // Check for a collision
 
186
                                if( dist < minDistance )
 
187
                                {
 
188
                                        // Project back to surface along normal
 
189
                                        position = position + (float4)((minDistance - dist)*normalVector*0.9f);
 
190
                                        velocity = (position - previousPosition) * velocityCoefficient * isolverdt;
 
191
                                        float4 relativeVelocity = velocity - velocityOfSurfacePoint;
 
192
 
 
193
                                        float4 p1 = (float4)(normalize(cross(normalVector, segment)).xyz, 0.f);
 
194
                                        float4 p2 = (float4)(normalize(cross(p1, normalVector)).xyz, 0.f);
 
195
                                        // Full friction is sum of velocities in each direction of plane
 
196
                                        float4 frictionVector = p1*mydot3(relativeVelocity, p1) + p2*mydot3(relativeVelocity, p2);
 
197
 
 
198
                                        // Real friction is peak friction corrected by friction coefficients
 
199
                                        frictionVector = frictionVector * (colliderFriction*clothFriction);
 
200
 
 
201
                                        float approachSpeed = dot(relativeVelocity, normalVector);
 
202
 
 
203
                                        if( approachSpeed <= 0.0f )
 
204
                                                forceOnVertex -= frictionVector;
 
205
                                }
 
206
                                        
 
207
                        }
 
208
                }
 
209
        }
 
210
        
 
211
        g_vertexVelocities[nodeID] = (float4)(velocity.xyz, 0.f);       
 
212
 
 
213
        // Update external force
 
214
        g_vertexForces[nodeID] = (float4)(forceOnVertex.xyz, 0.f);
 
215
 
 
216
        g_vertexPositions[nodeID] = (float4)(position.xyz, 0.f);
 
217
}
 
218
 
 
219
);