Bullet Collision Detection & Physics Library
btSoftBodySolver_OpenCLSIMDAware.cpp
Go to the documentation of this file.
1 /*
2 Bullet Continuous Collision Detection and Physics Library
3 Copyright (c) 2003-2006 Erwin Coumans http://continuousphysics.com/Bullet/
4 
5 This software is provided 'as-is', without any express or implied warranty.
6 In no event will the authors be held liable for any damages arising from the use of this software.
7 Permission is granted to anyone to use this software for any purpose,
8 including commercial applications, and to alter it and redistribute it freely,
9 subject to the following restrictions:
10 
11 1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required.
12 2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.
13 3. This notice may not be removed or altered from any source distribution.
14 */
15 
16 
17 #include "BulletCollision/CollisionShapes/btTriangleIndexVertexArray.h"
18 #include "vectormath/vmInclude.h"
19 #include <stdio.h> //@todo: remove the debugging printf at some stage
20 #include "btSoftBodySolver_OpenCLSIMDAware.h"
21 #include "BulletSoftBody/btSoftBodySolverVertexBuffer.h"
22 #include "BulletSoftBody/btSoftBody.h"
23 #include "BulletCollision/CollisionShapes/btCapsuleShape.h"
24 #include <limits.h>
25 
26 #define WAVEFRONT_SIZE 32
27 #define WAVEFRONT_BLOCK_MULTIPLIER 2
28 #define GROUP_SIZE (WAVEFRONT_SIZE*WAVEFRONT_BLOCK_MULTIPLIER)
29 #define LINKS_PER_SIMD_LANE 16
30 
31 static const size_t workGroupSize = GROUP_SIZE;
32 
33 
34 //CL_VERSION_1_1 seems broken on NVidia SDK so just disable it
35 
37 #define MSTRINGIFY(A) #A
38 static const char* UpdatePositionsFromVelocitiesCLString =
39 #include "OpenCLC10/UpdatePositionsFromVelocities.cl"
40 static const char* SolvePositionsCLString =
41 #include "OpenCLC10/SolvePositionsSIMDBatched.cl"
42 static const char* UpdateNodesCLString =
43 #include "OpenCLC10/UpdateNodes.cl"
44 static const char* UpdatePositionsCLString =
45 #include "OpenCLC10/UpdatePositions.cl"
46 static const char* UpdateConstantsCLString =
47 #include "OpenCLC10/UpdateConstants.cl"
48 static const char* IntegrateCLString =
49 #include "OpenCLC10/Integrate.cl"
50 static const char* ApplyForcesCLString =
51 #include "OpenCLC10/ApplyForces.cl"
52 static const char* UpdateFixedVertexPositionsCLString =
53 #include "OpenCLC10/UpdateFixedVertexPositions.cl"
54 static const char* UpdateNormalsCLString =
55 #include "OpenCLC10/UpdateNormals.cl"
56 static const char* VSolveLinksCLString =
57 #include "OpenCLC10/VSolveLinks.cl"
58 static const char* SolveCollisionsAndUpdateVelocitiesCLString =
59 #include "OpenCLC10/SolveCollisionsAndUpdateVelocitiesSIMDBatched.cl"
60 static const char* OutputToVertexArrayCLString =
61 #include "OpenCLC10/OutputToVertexArray.cl"
62 
63 
64 
65 btSoftBodyLinkDataOpenCLSIMDAware::btSoftBodyLinkDataOpenCLSIMDAware(cl_command_queue queue, cl_context ctx) :
66  m_cqCommandQue(queue),
67  m_wavefrontSize( WAVEFRONT_SIZE ),
68  m_linksPerWorkItem( LINKS_PER_SIMD_LANE ),
69  m_maxBatchesWithinWave( 0 ),
70  m_maxLinksPerWavefront( m_wavefrontSize * m_linksPerWorkItem ),
71  m_numWavefronts( 0 ),
72  m_maxVertex( 0 ),
73  m_clNumBatchesAndVerticesWithinWaves( queue, ctx, &m_numBatchesAndVerticesWithinWaves, true ),
74  m_clWavefrontVerticesGlobalAddresses( queue, ctx, &m_wavefrontVerticesGlobalAddresses, true ),
75  m_clLinkVerticesLocalAddresses( queue, ctx, &m_linkVerticesLocalAddresses, true ),
76  m_clLinkStrength( queue, ctx, &m_linkStrength, false ),
77  m_clLinksMassLSC( queue, ctx, &m_linksMassLSC, false ),
78  m_clLinksRestLengthSquared( queue, ctx, &m_linksRestLengthSquared, false ),
79  m_clLinksRestLength( queue, ctx, &m_linksRestLength, false ),
80  m_clLinksMaterialLinearStiffnessCoefficient( queue, ctx, &m_linksMaterialLinearStiffnessCoefficient, false )
81 {
82 }
83 
84 btSoftBodyLinkDataOpenCLSIMDAware::~btSoftBodyLinkDataOpenCLSIMDAware()
85 {
86 }
87 
88 static Vectormath::Aos::Vector3 toVector3( const btVector3 &vec )
89 {
90  Vectormath::Aos::Vector3 outVec( vec.getX(), vec.getY(), vec.getZ() );
91  return outVec;
92 }
93 
95 void btSoftBodyLinkDataOpenCLSIMDAware::createLinks( int numLinks )
96 {
97  int previousSize = m_links.size();
98  int newSize = previousSize + numLinks;
99 
100  btSoftBodyLinkData::createLinks( numLinks );
101 
102  // Resize the link addresses array as well
103  m_linkAddresses.resize( newSize );
104 }
105 
107 void btSoftBodyLinkDataOpenCLSIMDAware::setLinkAt(
108  const LinkDescription &link,
109  int linkIndex )
110 {
111  btSoftBodyLinkData::setLinkAt( link, linkIndex );
112 
113  if( link.getVertex0() > m_maxVertex )
114  m_maxVertex = link.getVertex0();
115  if( link.getVertex1() > m_maxVertex )
116  m_maxVertex = link.getVertex1();
117 
118  // Set the link index correctly for initialisation
119  m_linkAddresses[linkIndex] = linkIndex;
120 }
121 
122 bool btSoftBodyLinkDataOpenCLSIMDAware::onAccelerator()
123 {
124  return m_onGPU;
125 }
126 
127 bool btSoftBodyLinkDataOpenCLSIMDAware::moveToAccelerator()
128 {
129  bool success = true;
130  success = success && m_clNumBatchesAndVerticesWithinWaves.moveToGPU();
131  success = success && m_clWavefrontVerticesGlobalAddresses.moveToGPU();
132  success = success && m_clLinkVerticesLocalAddresses.moveToGPU();
133  success = success && m_clLinkStrength.moveToGPU();
134  success = success && m_clLinksMassLSC.moveToGPU();
135  success = success && m_clLinksRestLengthSquared.moveToGPU();
136  success = success && m_clLinksRestLength.moveToGPU();
137  success = success && m_clLinksMaterialLinearStiffnessCoefficient.moveToGPU();
138 
139  if( success ) {
140  m_onGPU = true;
141  }
142 
143  return success;
144 }
145 
146 bool btSoftBodyLinkDataOpenCLSIMDAware::moveFromAccelerator()
147 {
148  bool success = true;
149  success = success && m_clNumBatchesAndVerticesWithinWaves.moveToGPU();
150  success = success && m_clWavefrontVerticesGlobalAddresses.moveToGPU();
151  success = success && m_clLinkVerticesLocalAddresses.moveToGPU();
152  success = success && m_clLinkStrength.moveFromGPU();
153  success = success && m_clLinksMassLSC.moveFromGPU();
154  success = success && m_clLinksRestLengthSquared.moveFromGPU();
155  success = success && m_clLinksRestLength.moveFromGPU();
156  success = success && m_clLinksMaterialLinearStiffnessCoefficient.moveFromGPU();
157 
158  if( success ) {
159  m_onGPU = false;
160  }
161 
162  return success;
163 }
164 
165 
166 
167 
168 
169 
170 
171 
172 btOpenCLSoftBodySolverSIMDAware::btOpenCLSoftBodySolverSIMDAware(cl_command_queue queue, cl_context ctx, bool bUpdateAchchoredNodePos) :
173  btOpenCLSoftBodySolver( queue, ctx, bUpdateAchchoredNodePos ),
174  m_linkData(queue, ctx)
175 {
176  // Initial we will clearly need to update solver constants
177  // For now this is global for the cloths linked with this solver - we should probably make this body specific
178  // for performance in future once we understand more clearly when constants need to be updated
179  m_updateSolverConstants = true;
180 
181  m_shadersInitialized = false;
182 }
183 
184 btOpenCLSoftBodySolverSIMDAware::~btOpenCLSoftBodySolverSIMDAware()
185 {
186  releaseKernels();
187 }
188 
189 void btOpenCLSoftBodySolverSIMDAware::optimize( btAlignedObjectArray< btSoftBody * > &softBodies ,bool forceUpdate)
190 {
191  if( forceUpdate || m_softBodySet.size() != softBodies.size() )
192  {
193  // Have a change in the soft body set so update, reloading all the data
194  getVertexData().clear();
195  getTriangleData().clear();
196  getLinkData().clear();
197  m_softBodySet.resize(0);
198  m_anchorIndex.clear();
199 
200  int maxPiterations = 0;
201  int maxViterations = 0;
202 
203  for( int softBodyIndex = 0; softBodyIndex < softBodies.size(); ++softBodyIndex )
204  {
205  btSoftBody *softBody = softBodies[ softBodyIndex ];
206  using Vectormath::Aos::Matrix3;
207  using Vectormath::Aos::Point3;
208 
209  // Create SoftBody that will store the information within the solver
210  btOpenCLAcceleratedSoftBodyInterface* newSoftBody = new btOpenCLAcceleratedSoftBodyInterface( softBody );
211  m_softBodySet.push_back( newSoftBody );
212 
213  m_perClothAcceleration.push_back( toVector3(softBody->getWorldInfo()->m_gravity) );
214  m_perClothDampingFactor.push_back(softBody->m_cfg.kDP);
215  m_perClothVelocityCorrectionCoefficient.push_back( softBody->m_cfg.kVCF );
216  m_perClothLiftFactor.push_back( softBody->m_cfg.kLF );
217  m_perClothDragFactor.push_back( softBody->m_cfg.kDG );
218  m_perClothMediumDensity.push_back(softBody->getWorldInfo()->air_density);
219  // Simple init values. Actually we'll put 0 and -1 into them at the appropriate time
220  m_perClothFriction.push_back(softBody->m_cfg.kDF);
221  m_perClothCollisionObjects.push_back( CollisionObjectIndices(-1, -1) );
222 
223  // Add space for new vertices and triangles in the default solver for now
224  // TODO: Include space here for tearing too later
225  int firstVertex = getVertexData().getNumVertices();
226  int numVertices = softBody->m_nodes.size();
227  // Round maxVertices to a multiple of the workgroup size so we know we're safe to run over in a given group
228  // maxVertices can be increased to allow tearing, but should be used sparingly because these extra verts will always be processed
229  int maxVertices = GROUP_SIZE*((numVertices+GROUP_SIZE)/GROUP_SIZE);
230  // Allocate space for new vertices in all the vertex arrays
231  getVertexData().createVertices( numVertices, softBodyIndex, maxVertices );
232 
233 
234  int firstTriangle = getTriangleData().getNumTriangles();
235  int numTriangles = softBody->m_faces.size();
236  int maxTriangles = numTriangles;
237  getTriangleData().createTriangles( maxTriangles );
238 
239  // Copy vertices from softbody into the solver
240  for( int vertex = 0; vertex < numVertices; ++vertex )
241  {
242  Point3 multPoint(softBody->m_nodes[vertex].m_x.getX(), softBody->m_nodes[vertex].m_x.getY(), softBody->m_nodes[vertex].m_x.getZ());
243  btSoftBodyVertexData::VertexDescription desc;
244 
245  // TODO: Position in the softbody might be pre-transformed
246  // or we may need to adapt for the pose.
247  //desc.setPosition( cloth.getMeshTransform()*multPoint );
248  desc.setPosition( multPoint );
249 
250  float vertexInverseMass = softBody->m_nodes[vertex].m_im;
251  desc.setInverseMass(vertexInverseMass);
252  getVertexData().setVertexAt( desc, firstVertex + vertex );
253 
254  m_anchorIndex.push_back(-1);
255  }
256  for( int vertex = numVertices; vertex < maxVertices; ++vertex )
257  {
258  m_anchorIndex.push_back(-1.0);
259  }
260 
261  // Copy triangles similarly
262  // We're assuming here that vertex indices are based on the firstVertex rather than the entire scene
263  for( int triangle = 0; triangle < numTriangles; ++triangle )
264  {
265  // Note that large array storage is relative to the array not to the cloth
266  // So we need to add firstVertex to each value
267  int vertexIndex0 = (softBody->m_faces[triangle].m_n[0] - &(softBody->m_nodes[0]));
268  int vertexIndex1 = (softBody->m_faces[triangle].m_n[1] - &(softBody->m_nodes[0]));
269  int vertexIndex2 = (softBody->m_faces[triangle].m_n[2] - &(softBody->m_nodes[0]));
270  btSoftBodyTriangleData::TriangleDescription newTriangle(vertexIndex0 + firstVertex, vertexIndex1 + firstVertex, vertexIndex2 + firstVertex);
271  getTriangleData().setTriangleAt( newTriangle, firstTriangle + triangle );
272 
273  // Increase vertex triangle counts for this triangle
274  getVertexData().getTriangleCount(newTriangle.getVertexSet().vertex0)++;
275  getVertexData().getTriangleCount(newTriangle.getVertexSet().vertex1)++;
276  getVertexData().getTriangleCount(newTriangle.getVertexSet().vertex2)++;
277  }
278 
279  int firstLink = getLinkData().getNumLinks();
280  int numLinks = softBody->m_links.size();
281  int maxLinks = numLinks;
282 
283  // Allocate space for the links
284  getLinkData().createLinks( numLinks );
285 
286  // Add the links
287  for( int link = 0; link < numLinks; ++link )
288  {
289  int vertexIndex0 = softBody->m_links[link].m_n[0] - &(softBody->m_nodes[0]);
290  int vertexIndex1 = softBody->m_links[link].m_n[1] - &(softBody->m_nodes[0]);
291 
292  btSoftBodyLinkData::LinkDescription newLink(vertexIndex0 + firstVertex, vertexIndex1 + firstVertex, softBody->m_links[link].m_material->m_kLST);
293  newLink.setLinkStrength(1.f);
294  getLinkData().setLinkAt(newLink, firstLink + link);
295  }
296 
297  newSoftBody->setFirstVertex( firstVertex );
298  newSoftBody->setFirstTriangle( firstTriangle );
299  newSoftBody->setNumVertices( numVertices );
300  newSoftBody->setMaxVertices( maxVertices );
301  newSoftBody->setNumTriangles( numTriangles );
302  newSoftBody->setMaxTriangles( maxTriangles );
303  newSoftBody->setFirstLink( firstLink );
304  newSoftBody->setNumLinks( numLinks );
305 
306  // Find maximum piterations and viterations
307  int piterations = softBody->m_cfg.piterations;
308 
309  if ( piterations > maxPiterations )
310  maxPiterations = piterations;
311 
312  int viterations = softBody->m_cfg.viterations;
313 
314  if ( viterations > maxViterations )
315  maxViterations = viterations;
316 
317  // zero mass
318  for( int vertex = 0; vertex < numVertices; ++vertex )
319  {
320  if ( softBody->m_nodes[vertex].m_im == 0 )
321  {
322  AnchorNodeInfoCL nodeInfo;
323  nodeInfo.clVertexIndex = firstVertex + vertex;
324  nodeInfo.pNode = &softBody->m_nodes[vertex];
325 
326  m_anchorNodeInfoArray.push_back(nodeInfo);
327  }
328  }
329 
330  // anchor position
331  if ( numVertices > 0 )
332  {
333  for ( int anchorIndex = 0; anchorIndex < softBody->m_anchors.size(); anchorIndex++ )
334  {
335  btSoftBody::Node* anchorNode = softBody->m_anchors[anchorIndex].m_node;
336  btSoftBody::Node* firstNode = &softBody->m_nodes[0];
337 
338  AnchorNodeInfoCL nodeInfo;
339  nodeInfo.clVertexIndex = firstVertex + (int)(anchorNode - firstNode);
340  nodeInfo.pNode = anchorNode;
341 
342  m_anchorNodeInfoArray.push_back(nodeInfo);
343  }
344  }
345  }
346 
347  m_anchorPosition.clear();
348  m_anchorPosition.resize(m_anchorNodeInfoArray.size());
349 
350  for ( int anchorNode = 0; anchorNode < m_anchorNodeInfoArray.size(); anchorNode++ )
351  {
352  const AnchorNodeInfoCL& anchorNodeInfo = m_anchorNodeInfoArray[anchorNode];
353  m_anchorIndex[anchorNodeInfo.clVertexIndex] = anchorNode;
354  getVertexData().getInverseMass(anchorNodeInfo.clVertexIndex) = 0.0f;
355  }
356 
357  updateConstants(0.f);
358 
359  // set position and velocity iterations
360  setNumberOfPositionIterations(maxPiterations);
361  setNumberOfVelocityIterations(maxViterations);
362 
363  // set wind velocity
364  m_perClothWindVelocity.resize( m_softBodySet.size() );
365  for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex )
366  {
367  btSoftBody *softBody = m_softBodySet[softBodyIndex]->getSoftBody();
368  m_perClothWindVelocity[softBodyIndex] = toVector3(softBody->getWindVelocity());
369  }
370 
371  m_clPerClothWindVelocity.changedOnCPU();
372 
373  // generate batches
374  m_linkData.generateBatches();
375  m_triangleData.generateBatches();
376 
377  // Build the shaders to match the batching parameters
378  buildShaders();
379  }
380 }
381 
382 
383 btSoftBodyLinkData &btOpenCLSoftBodySolverSIMDAware::getLinkData()
384 {
385  // TODO: Consider setting link data to "changed" here
386  return m_linkData;
387 }
388 
389 
390 
391 
392 void btOpenCLSoftBodySolverSIMDAware::updateConstants( float timeStep )
393 {
394 
395  using namespace Vectormath::Aos;
396 
397  if( m_updateSolverConstants )
398  {
399  m_updateSolverConstants = false;
400 
401  // Will have to redo this if we change the structure (tear, maybe) or various other possible changes
402 
403  // Initialise link constants
404  const int numLinks = m_linkData.getNumLinks();
405  for( int linkIndex = 0; linkIndex < numLinks; ++linkIndex )
406  {
407  btSoftBodyLinkData::LinkNodePair &vertices( m_linkData.getVertexPair(linkIndex) );
408  m_linkData.getRestLength(linkIndex) = length((m_vertexData.getPosition( vertices.vertex0 ) - m_vertexData.getPosition( vertices.vertex1 )));
409  float invMass0 = m_vertexData.getInverseMass(vertices.vertex0);
410  float invMass1 = m_vertexData.getInverseMass(vertices.vertex1);
411  float linearStiffness = m_linkData.getLinearStiffnessCoefficient(linkIndex);
412  float massLSC = (invMass0 + invMass1)/linearStiffness;
413  m_linkData.getMassLSC(linkIndex) = massLSC;
414  float restLength = m_linkData.getRestLength(linkIndex);
415  float restLengthSquared = restLength*restLength;
416  m_linkData.getRestLengthSquared(linkIndex) = restLengthSquared;
417  }
418  }
419 
420 }
421 
422 
423 
424 void btOpenCLSoftBodySolverSIMDAware::solveConstraints( float solverdt )
425 {
426 
427  using Vectormath::Aos::Vector3;
428  using Vectormath::Aos::Point3;
429  using Vectormath::Aos::lengthSqr;
430  using Vectormath::Aos::dot;
431 
432  // Prepare links
433  int numLinks = m_linkData.getNumLinks();
434  int numVertices = m_vertexData.getNumVertices();
435 
436  float kst = 1.f;
437  float ti = 0.f;
438 
439 
440  m_clPerClothDampingFactor.moveToGPU();
441  m_clPerClothVelocityCorrectionCoefficient.moveToGPU();
442 
443 
444  // Ensure data is on accelerator
445  m_linkData.moveToAccelerator();
446  m_vertexData.moveToAccelerator();
447 
448 
449  //prepareLinks();
450 
451  prepareCollisionConstraints();
452 
453  // Solve drift
454  for( int iteration = 0; iteration < m_numberOfPositionIterations ; ++iteration )
455  {
456 
457  for( int i = 0; i < m_linkData.m_wavefrontBatchStartLengths.size(); ++i )
458  {
459  int startWave = m_linkData.m_wavefrontBatchStartLengths[i].start;
460  int numWaves = m_linkData.m_wavefrontBatchStartLengths[i].length;
461  solveLinksForPosition( startWave, numWaves, kst, ti );
462  }
463  } // for( int iteration = 0; iteration < m_numberOfPositionIterations ; ++iteration )
464 
465 
466  // At this point assume that the force array is blank - we will overwrite it
467  solveCollisionsAndUpdateVelocities( 1.f/solverdt );
468 }
469 
470 
472 // Kernel dispatches
473 
474 
475 void btOpenCLSoftBodySolverSIMDAware::solveLinksForPosition( int startWave, int numWaves, float kst, float ti )
476 {
477  cl_int ciErrNum;
478  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,0, sizeof(int), &startWave);
479  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,1, sizeof(int), &numWaves);
480  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,2, sizeof(float), &kst);
481  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,3, sizeof(float), &ti);
482 
483 
484  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,4, sizeof(cl_mem), &m_linkData.m_clNumBatchesAndVerticesWithinWaves.m_buffer);
485  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,5, sizeof(cl_mem), &m_linkData.m_clWavefrontVerticesGlobalAddresses.m_buffer);
486  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,6, sizeof(cl_mem), &m_linkData.m_clLinkVerticesLocalAddresses.m_buffer);
487  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,7, sizeof(cl_mem), &m_linkData.m_clLinksMassLSC.m_buffer);
488 
489  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,8, sizeof(cl_mem), &m_linkData.m_clLinksRestLengthSquared.m_buffer);
490  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,9, sizeof(cl_mem), &m_vertexData.m_clVertexInverseMass.m_buffer);
491  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,10, sizeof(cl_mem), &m_vertexData.m_clVertexPosition.m_buffer);
492 
493  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,11, WAVEFRONT_BLOCK_MULTIPLIER*sizeof(cl_int2), 0);
494  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,12, m_linkData.getMaxVerticesPerWavefront()*WAVEFRONT_BLOCK_MULTIPLIER*sizeof(cl_float4), 0);
495  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,13, m_linkData.getMaxVerticesPerWavefront()*WAVEFRONT_BLOCK_MULTIPLIER*sizeof(cl_float), 0);
496 
497  size_t numWorkItems = workGroupSize*((numWaves*WAVEFRONT_SIZE + (workGroupSize-1)) / workGroupSize);
498 
499  ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,m_solvePositionsFromLinksKernel,1,NULL,&numWorkItems,&workGroupSize,0,0,0);
500 
501  if( ciErrNum!= CL_SUCCESS )
502  {
503  btAssert( 0 && "enqueueNDRangeKernel(m_solvePositionsFromLinksKernel)");
504  }
505 
506 } // solveLinksForPosition
507 
508 void btOpenCLSoftBodySolverSIMDAware::solveCollisionsAndUpdateVelocities( float isolverdt )
509 {
510  // Copy kernel parameters to GPU
511  m_vertexData.moveToAccelerator();
512  m_clPerClothFriction.moveToGPU();
513  m_clPerClothDampingFactor.moveToGPU();
514  m_clPerClothCollisionObjects.moveToGPU();
515  m_clCollisionObjectDetails.moveToGPU();
516 
517  cl_int ciErrNum;
518  int numVerts = m_vertexData.getNumVertices();
519  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 0, sizeof(int), &numVerts);
520  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 1, sizeof(int), &isolverdt);
521  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 2, sizeof(cl_mem),&m_vertexData.m_clClothIdentifier.m_buffer);
522  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 3, sizeof(cl_mem),&m_vertexData.m_clVertexPreviousPosition.m_buffer);
523  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 4, sizeof(cl_mem),&m_clPerClothFriction.m_buffer);
524  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 5, sizeof(cl_mem),&m_clPerClothDampingFactor.m_buffer);
525  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 6, sizeof(cl_mem),&m_clPerClothCollisionObjects.m_buffer);
526  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 7, sizeof(cl_mem),&m_clCollisionObjectDetails.m_buffer);
527  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 8, sizeof(cl_mem),&m_vertexData.m_clVertexForceAccumulator.m_buffer);
528  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 9, sizeof(cl_mem),&m_vertexData.m_clVertexVelocity.m_buffer);
529  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 10, sizeof(cl_mem),&m_vertexData.m_clVertexPosition.m_buffer);
530  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 11, sizeof(CollisionShapeDescription)*16,0);
531  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 12, sizeof(cl_mem),&m_vertexData.m_clVertexInverseMass.m_buffer);
532  size_t numWorkItems = workGroupSize*((m_vertexData.getNumVertices() + (workGroupSize-1)) / workGroupSize);
533 
534  if (numWorkItems)
535  {
536  ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,m_solveCollisionsAndUpdateVelocitiesKernel, 1, NULL, &numWorkItems, &workGroupSize,0,0,0);
537 
538  if( ciErrNum != CL_SUCCESS )
539  {
540  btAssert( 0 && "enqueueNDRangeKernel(m_solveCollisionsAndUpdateVelocitiesKernel)");
541  }
542  }
543 
544 } // btOpenCLSoftBodySolverSIMDAware::updateVelocitiesFromPositionsWithoutVelocities
545 
546 // End kernel dispatches
548 
549 
550 
551 bool btOpenCLSoftBodySolverSIMDAware::buildShaders()
552 {
553  releaseKernels();
554 
555  if( m_shadersInitialized )
556  return true;
557 
558  const char* additionalMacros="";
559 
560  m_currentCLFunctions->clearKernelCompilationFailures();
561 
562  char *wavefrontMacros = new char[256];
563 
564  sprintf(
565  wavefrontMacros,
566  "-DMAX_NUM_VERTICES_PER_WAVE=%d -DMAX_BATCHES_PER_WAVE=%d -DWAVEFRONT_SIZE=%d -DWAVEFRONT_BLOCK_MULTIPLIER=%d -DBLOCK_SIZE=%d",
567  m_linkData.getMaxVerticesPerWavefront(),
568  m_linkData.getMaxBatchesPerWavefront(),
569  m_linkData.getWavefrontSize(),
570  WAVEFRONT_BLOCK_MULTIPLIER,
571  WAVEFRONT_BLOCK_MULTIPLIER*m_linkData.getWavefrontSize());
572 
573  m_updatePositionsFromVelocitiesKernel = m_currentCLFunctions->compileCLKernelFromString( UpdatePositionsFromVelocitiesCLString, "UpdatePositionsFromVelocitiesKernel", additionalMacros,"OpenCLC10/UpdatePositionsFromVelocities.cl");
574  m_solvePositionsFromLinksKernel = m_currentCLFunctions->compileCLKernelFromString( SolvePositionsCLString, "SolvePositionsFromLinksKernel", wavefrontMacros ,"OpenCLC10/SolvePositionsSIMDBatched.cl");
575  m_updateVelocitiesFromPositionsWithVelocitiesKernel = m_currentCLFunctions->compileCLKernelFromString( UpdateNodesCLString, "updateVelocitiesFromPositionsWithVelocitiesKernel", additionalMacros ,"OpenCLC10/UpdateNodes.cl");
576  m_updateVelocitiesFromPositionsWithoutVelocitiesKernel = m_currentCLFunctions->compileCLKernelFromString( UpdatePositionsCLString, "updateVelocitiesFromPositionsWithoutVelocitiesKernel", additionalMacros,"OpenCLC10/UpdatePositions.cl");
577  m_integrateKernel = m_currentCLFunctions->compileCLKernelFromString( IntegrateCLString, "IntegrateKernel", additionalMacros ,"OpenCLC10/Integrate.cl");
578  m_applyForcesKernel = m_currentCLFunctions->compileCLKernelFromString( ApplyForcesCLString, "ApplyForcesKernel", additionalMacros,"OpenCLC10/ApplyForces.cl" );
579  m_updateFixedVertexPositionsKernel = m_currentCLFunctions->compileCLKernelFromString( UpdateFixedVertexPositionsCLString, "UpdateFixedVertexPositions" ,additionalMacros,"OpenCLC10/UpdateFixedVertexPositions.cl");
580  m_solveCollisionsAndUpdateVelocitiesKernel = m_currentCLFunctions->compileCLKernelFromString( SolveCollisionsAndUpdateVelocitiesCLString, "SolveCollisionsAndUpdateVelocitiesKernel", additionalMacros ,"OpenCLC10/SolveCollisionsAndUpdateVelocitiesSIMDBatched.cl");
581 
582  // TODO: Rename to UpdateSoftBodies
583  m_resetNormalsAndAreasKernel = m_currentCLFunctions->compileCLKernelFromString( UpdateNormalsCLString, "ResetNormalsAndAreasKernel", additionalMacros ,"OpenCLC10/UpdateNormals.cl");
584  m_normalizeNormalsAndAreasKernel = m_currentCLFunctions->compileCLKernelFromString( UpdateNormalsCLString, "NormalizeNormalsAndAreasKernel", additionalMacros ,"OpenCLC10/UpdateNormals.cl");
585  m_updateSoftBodiesKernel = m_currentCLFunctions->compileCLKernelFromString( UpdateNormalsCLString, "UpdateSoftBodiesKernel", additionalMacros ,"OpenCLC10/UpdateNormals.cl");
586 
587  delete [] wavefrontMacros;
588 
589  if( m_currentCLFunctions->getKernelCompilationFailures()==0)
590  {
591  m_shadersInitialized = true;
592  }
593 
594  return m_shadersInitialized;
595 }
596 
597 
598 
599 
600 static Vectormath::Aos::Transform3 toTransform3( const btTransform &transform )
601 {
602  Vectormath::Aos::Transform3 outTransform;
603  outTransform.setCol(0, toVector3(transform.getBasis().getColumn(0)));
604  outTransform.setCol(1, toVector3(transform.getBasis().getColumn(1)));
605  outTransform.setCol(2, toVector3(transform.getBasis().getColumn(2)));
606  outTransform.setCol(3, toVector3(transform.getOrigin()));
607  return outTransform;
608 }
609 
610 
611 static void generateBatchesOfWavefronts( btAlignedObjectArray < btAlignedObjectArray <int> > &linksForWavefronts, btSoftBodyLinkData &linkData, int numVertices, btAlignedObjectArray < btAlignedObjectArray <int> > &wavefrontBatches )
612 {
613  // A per-batch map of truth values stating whether a given vertex is in that batch
614  // This allows us to significantly optimize the batching
615  btAlignedObjectArray <btAlignedObjectArray<bool> > mapOfVerticesInBatches;
616 
617  for( int waveIndex = 0; waveIndex < linksForWavefronts.size(); ++waveIndex )
618  {
619  btAlignedObjectArray <int> &wavefront( linksForWavefronts[waveIndex] );
620 
621  int batch = 0;
622  bool placed = false;
623  while( batch < wavefrontBatches.size() && !placed )
624  {
625  // Test the current batch, see if this wave shares any vertex with the waves in the batch
626  bool foundSharedVertex = false;
627  for( int link = 0; link < wavefront.size(); ++link )
628  {
629  btSoftBodyLinkData::LinkNodePair vertices = linkData.getVertexPair( wavefront[link] );
630  if( (mapOfVerticesInBatches[batch])[vertices.vertex0] || (mapOfVerticesInBatches[batch])[vertices.vertex1] )
631  {
632  foundSharedVertex = true;
633  }
634  }
635 
636  if( !foundSharedVertex )
637  {
638  wavefrontBatches[batch].push_back( waveIndex );
639  // Insert vertices into this batch too
640  for( int link = 0; link < wavefront.size(); ++link )
641  {
642  btSoftBodyLinkData::LinkNodePair vertices = linkData.getVertexPair( wavefront[link] );
643  (mapOfVerticesInBatches[batch])[vertices.vertex0] = true;
644  (mapOfVerticesInBatches[batch])[vertices.vertex1] = true;
645  }
646  placed = true;
647  }
648  batch++;
649  }
650  if( batch == wavefrontBatches.size() && !placed )
651  {
652  wavefrontBatches.resize( batch + 1 );
653  wavefrontBatches[batch].push_back( waveIndex );
654 
655  // And resize map as well
656  mapOfVerticesInBatches.resize( batch + 1 );
657 
658  // Resize maps with total number of vertices
659  mapOfVerticesInBatches[batch].resize( numVertices+1, false );
660 
661  // Insert vertices into this batch too
662  for( int link = 0; link < wavefront.size(); ++link )
663  {
664  btSoftBodyLinkData::LinkNodePair vertices = linkData.getVertexPair( wavefront[link] );
665  (mapOfVerticesInBatches[batch])[vertices.vertex0] = true;
666  (mapOfVerticesInBatches[batch])[vertices.vertex1] = true;
667  }
668  }
669  }
670  mapOfVerticesInBatches.clear();
671 }
672 
673 // Function to remove an object from a vector maintaining correct ordering of the vector
674 template< typename T > static void removeFromVector( btAlignedObjectArray< T > &vectorToUpdate, int indexToRemove )
675 {
676  int currentSize = vectorToUpdate.size();
677  for( int i = indexToRemove; i < (currentSize-1); ++i )
678  {
679  vectorToUpdate[i] = vectorToUpdate[i+1];
680  }
681  if( currentSize > 0 )
682  vectorToUpdate.resize( currentSize - 1 );
683 }
684 
688 template< typename T > static void insertAtIndex( btAlignedObjectArray< T > &vectorToUpdate, int index, T element )
689 {
690  vectorToUpdate.resize( vectorToUpdate.size() + 1 );
691  for( int i = (vectorToUpdate.size() - 1); i > index; --i )
692  {
693  vectorToUpdate[i] = vectorToUpdate[i-1];
694  }
695  vectorToUpdate[index] = element;
696 }
697 
702 template< typename T > static void insertUniqueAndOrderedIntoVector( btAlignedObjectArray<T> &vectorToUpdate, T element )
703 {
704  int index = 0;
705  while( index < vectorToUpdate.size() && vectorToUpdate[index] < element )
706  {
707  index++;
708  }
709  if( index == vectorToUpdate.size() || vectorToUpdate[index] != element )
710  insertAtIndex( vectorToUpdate, index, element );
711 }
712 
713 static void generateLinksPerVertex( int numVertices, btSoftBodyLinkData &linkData, btAlignedObjectArray< int > &listOfLinksPerVertex, btAlignedObjectArray <int> &numLinksPerVertex, int &maxLinks )
714 {
715  for( int linkIndex = 0; linkIndex < linkData.getNumLinks(); ++linkIndex )
716  {
717  btSoftBodyLinkData::LinkNodePair nodes( linkData.getVertexPair(linkIndex) );
718  numLinksPerVertex[nodes.vertex0]++;
719  numLinksPerVertex[nodes.vertex1]++;
720  }
721  int maxLinksPerVertex = 0;
722  for( int vertexIndex = 0; vertexIndex < numVertices; ++vertexIndex )
723  {
724  maxLinksPerVertex = btMax(numLinksPerVertex[vertexIndex], maxLinksPerVertex);
725  }
726  maxLinks = maxLinksPerVertex;
727 
728  btAlignedObjectArray< int > linksFoundPerVertex;
729  linksFoundPerVertex.resize( numVertices, 0 );
730 
731  listOfLinksPerVertex.resize( maxLinksPerVertex * numVertices );
732 
733  for( int linkIndex = 0; linkIndex < linkData.getNumLinks(); ++linkIndex )
734  {
735  btSoftBodyLinkData::LinkNodePair nodes( linkData.getVertexPair(linkIndex) );
736  {
737  // Do vertex 0
738  int vertexIndex = nodes.vertex0;
739  int linkForVertex = linksFoundPerVertex[nodes.vertex0];
740  int linkAddress = vertexIndex * maxLinksPerVertex + linkForVertex;
741 
742  listOfLinksPerVertex[linkAddress] = linkIndex;
743 
744  linksFoundPerVertex[nodes.vertex0] = linkForVertex + 1;
745  }
746  {
747  // Do vertex 1
748  int vertexIndex = nodes.vertex1;
749  int linkForVertex = linksFoundPerVertex[nodes.vertex1];
750  int linkAddress = vertexIndex * maxLinksPerVertex + linkForVertex;
751 
752  listOfLinksPerVertex[linkAddress] = linkIndex;
753 
754  linksFoundPerVertex[nodes.vertex1] = linkForVertex + 1;
755  }
756  }
757 }
758 
759 static void computeBatchingIntoWavefronts(
760  btSoftBodyLinkData &linkData,
761  int wavefrontSize,
762  int linksPerWorkItem,
763  int maxLinksPerWavefront,
764  btAlignedObjectArray < btAlignedObjectArray <int> > &linksForWavefronts,
765  btAlignedObjectArray< btAlignedObjectArray < btAlignedObjectArray <int> > > &batchesWithinWaves, /* wave, batch, links in batch */
766  btAlignedObjectArray< btAlignedObjectArray< int > > &verticesForWavefronts /* wavefront, vertex */
767  )
768 {
769 
770 
771  // Attempt generation of larger batches of links.
772  btAlignedObjectArray< bool > processedLink;
773  processedLink.resize( linkData.getNumLinks() );
774  btAlignedObjectArray< int > listOfLinksPerVertex;
775  int maxLinksPerVertex = 0;
776 
777  // Count num vertices
778  int numVertices = 0;
779  for( int linkIndex = 0; linkIndex < linkData.getNumLinks(); ++linkIndex )
780  {
781  btSoftBodyLinkData::LinkNodePair nodes( linkData.getVertexPair(linkIndex) );
782  numVertices = btMax( numVertices, nodes.vertex0 + 1 );
783  numVertices = btMax( numVertices, nodes.vertex1 + 1 );
784  }
785 
786  // Need list of links per vertex
787  // Compute valence of each vertex
788  btAlignedObjectArray <int> numLinksPerVertex;
789  numLinksPerVertex.resize(0);
790  numLinksPerVertex.resize( numVertices, 0 );
791 
792  generateLinksPerVertex( numVertices, linkData, listOfLinksPerVertex, numLinksPerVertex, maxLinksPerVertex );
793 
794  if (!numVertices)
795  return;
796 
797  for( int vertex = 0; vertex < 10; ++vertex )
798  {
799  for( int link = 0; link < numLinksPerVertex[vertex]; ++link )
800  {
801  int linkAddress = vertex * maxLinksPerVertex + link;
802  }
803  }
804 
805 
806  // At this point we know what links we have for each vertex so we can start batching
807 
808  // We want a vertex to start with, let's go with 0
809  int currentVertex = 0;
810  int linksProcessed = 0;
811 
812  btAlignedObjectArray <int> verticesToProcess;
813 
814  while( linksProcessed < linkData.getNumLinks() )
815  {
816  // Next wavefront
817  int nextWavefront = linksForWavefronts.size();
818  linksForWavefronts.resize( nextWavefront + 1 );
819  btAlignedObjectArray <int> &linksForWavefront(linksForWavefronts[nextWavefront]);
820  verticesForWavefronts.resize( nextWavefront + 1 );
821  btAlignedObjectArray<int> &vertexSet( verticesForWavefronts[nextWavefront] );
822 
823  linksForWavefront.resize(0);
824 
825  // Loop to find enough links to fill the wavefront
826  // Stopping if we either run out of links, or fill it
827  while( linksProcessed < linkData.getNumLinks() && linksForWavefront.size() < maxLinksPerWavefront )
828  {
829  // Go through the links for the current vertex
830  for( int link = 0; link < numLinksPerVertex[currentVertex] && linksForWavefront.size() < maxLinksPerWavefront; ++link )
831  {
832  int linkAddress = currentVertex * maxLinksPerVertex + link;
833  int linkIndex = listOfLinksPerVertex[linkAddress];
834 
835  // If we have not already processed this link, add it to the wavefront
836  // Claim it as another processed link
837  // Add the vertex at the far end to the list of vertices to process.
838  if( !processedLink[linkIndex] )
839  {
840  linksForWavefront.push_back( linkIndex );
841  linksProcessed++;
842  processedLink[linkIndex] = true;
843  int v0 = linkData.getVertexPair(linkIndex).vertex0;
844  int v1 = linkData.getVertexPair(linkIndex).vertex1;
845  if( v0 == currentVertex )
846  verticesToProcess.push_back( v1 );
847  else
848  verticesToProcess.push_back( v0 );
849  }
850  }
851  if( verticesToProcess.size() > 0 )
852  {
853  // Get the element on the front of the queue and remove it
854  currentVertex = verticesToProcess[0];
855  removeFromVector( verticesToProcess, 0 );
856  } else {
857  // If we've not yet processed all the links, find the first unprocessed one
858  // and select one of its vertices as the current vertex
859  if( linksProcessed < linkData.getNumLinks() )
860  {
861  int searchLink = 0;
862  while( processedLink[searchLink] )
863  searchLink++;
864  currentVertex = linkData.getVertexPair(searchLink).vertex0;
865  }
866  }
867  }
868 
869  // We have either finished or filled a wavefront
870  for( int link = 0; link < linksForWavefront.size(); ++link )
871  {
872  int v0 = linkData.getVertexPair( linksForWavefront[link] ).vertex0;
873  int v1 = linkData.getVertexPair( linksForWavefront[link] ).vertex1;
874  insertUniqueAndOrderedIntoVector( vertexSet, v0 );
875  insertUniqueAndOrderedIntoVector( vertexSet, v1 );
876  }
877  // Iterate over links mapped to the wave and batch those
878  // We can run a batch on each cycle trivially
879 
880  batchesWithinWaves.resize( batchesWithinWaves.size() + 1 );
881  btAlignedObjectArray < btAlignedObjectArray <int> > &batchesWithinWave( batchesWithinWaves[batchesWithinWaves.size()-1] );
882 
883 
884  for( int link = 0; link < linksForWavefront.size(); ++link )
885  {
886  int linkIndex = linksForWavefront[link];
887  btSoftBodyLinkData::LinkNodePair vertices = linkData.getVertexPair( linkIndex );
888 
889  int batch = 0;
890  bool placed = false;
891  while( batch < batchesWithinWave.size() && !placed )
892  {
893  bool foundSharedVertex = false;
894  if( batchesWithinWave[batch].size() >= wavefrontSize )
895  {
896  // If we have already filled this batch, move on to another
897  foundSharedVertex = true;
898  } else {
899  for( int link2 = 0; link2 < batchesWithinWave[batch].size(); ++link2 )
900  {
901  btSoftBodyLinkData::LinkNodePair vertices2 = linkData.getVertexPair( (batchesWithinWave[batch])[link2] );
902 
903  if( vertices.vertex0 == vertices2.vertex0 ||
904  vertices.vertex1 == vertices2.vertex0 ||
905  vertices.vertex0 == vertices2.vertex1 ||
906  vertices.vertex1 == vertices2.vertex1 )
907  {
908  foundSharedVertex = true;
909  break;
910  }
911  }
912  }
913  if( !foundSharedVertex )
914  {
915  batchesWithinWave[batch].push_back( linkIndex );
916  placed = true;
917  } else {
918  ++batch;
919  }
920  }
921  if( batch == batchesWithinWave.size() && !placed )
922  {
923  batchesWithinWave.resize( batch + 1 );
924  batchesWithinWave[batch].push_back( linkIndex );
925  }
926  }
927 
928  }
929 
930 }
931 
932 void btSoftBodyLinkDataOpenCLSIMDAware::generateBatches()
933 {
934  btAlignedObjectArray < btAlignedObjectArray <int> > linksForWavefronts;
935  btAlignedObjectArray < btAlignedObjectArray <int> > wavefrontBatches;
936  btAlignedObjectArray< btAlignedObjectArray < btAlignedObjectArray <int> > > batchesWithinWaves;
937  btAlignedObjectArray< btAlignedObjectArray< int > > verticesForWavefronts; // wavefronts, vertices in wavefront as an ordered set
938 
939  // Group the links into wavefronts
940  computeBatchingIntoWavefronts( *this, m_wavefrontSize, m_linksPerWorkItem, m_maxLinksPerWavefront, linksForWavefronts, batchesWithinWaves, verticesForWavefronts );
941 
942 
943  // Batch the wavefronts
944  generateBatchesOfWavefronts( linksForWavefronts, *this, m_maxVertex, wavefrontBatches );
945 
946  m_numWavefronts = linksForWavefronts.size();
947 
948  // At this point we have a description of which links we need to process in each wavefront
949 
950  // First correctly fill the batch ranges vector
951  int numBatches = wavefrontBatches.size();
952  m_wavefrontBatchStartLengths.resize(0);
953  int prefixSum = 0;
954  for( int batchIndex = 0; batchIndex < numBatches; ++batchIndex )
955  {
956  int wavesInBatch = wavefrontBatches[batchIndex].size();
957  int nextPrefixSum = prefixSum + wavesInBatch;
958  m_wavefrontBatchStartLengths.push_back( BatchPair( prefixSum, nextPrefixSum - prefixSum ) );
959 
960  prefixSum += wavesInBatch;
961  }
962 
963  // Also find max number of batches within a wave
964  m_maxBatchesWithinWave = 0;
965  m_maxVerticesWithinWave = 0;
966  m_numBatchesAndVerticesWithinWaves.resize( m_numWavefronts );
967  for( int waveIndex = 0; waveIndex < m_numWavefronts; ++waveIndex )
968  {
969  // See if the number of batches in this wave is greater than the current maxium
970  int batchesInCurrentWave = batchesWithinWaves[waveIndex].size();
971  int verticesInCurrentWave = verticesForWavefronts[waveIndex].size();
972  m_maxBatchesWithinWave = btMax( batchesInCurrentWave, m_maxBatchesWithinWave );
973  m_maxVerticesWithinWave = btMax( verticesInCurrentWave, m_maxVerticesWithinWave );
974  }
975 
976  // Add padding values both for alignment and as dudd addresses within LDS to compute junk rather than branch around
977  m_maxVerticesWithinWave = 16*((m_maxVerticesWithinWave/16)+2);
978 
979  // Now we know the maximum number of vertices per-wave we can resize the global vertices array
980  m_wavefrontVerticesGlobalAddresses.resize( m_maxVerticesWithinWave * m_numWavefronts );
981 
982  // Grab backup copies of all the link data arrays for the sorting process
983  btAlignedObjectArray<btSoftBodyLinkData::LinkNodePair> m_links_Backup(m_links);
984  btAlignedObjectArray<float> m_linkStrength_Backup(m_linkStrength);
985  btAlignedObjectArray<float> m_linksMassLSC_Backup(m_linksMassLSC);
986  btAlignedObjectArray<float> m_linksRestLengthSquared_Backup(m_linksRestLengthSquared);
987  //btAlignedObjectArray<Vectormath::Aos::Vector3> m_linksCLength_Backup(m_linksCLength);
988  //btAlignedObjectArray<float> m_linksLengthRatio_Backup(m_linksLengthRatio);
989  btAlignedObjectArray<float> m_linksRestLength_Backup(m_linksRestLength);
990  btAlignedObjectArray<float> m_linksMaterialLinearStiffnessCoefficient_Backup(m_linksMaterialLinearStiffnessCoefficient);
991 
992  // Resize to a wavefront sized batch per batch per wave so we get perfectly coherent memory accesses.
993  m_links.resize( m_maxBatchesWithinWave * m_wavefrontSize * m_numWavefronts );
994  m_linkVerticesLocalAddresses.resize( m_maxBatchesWithinWave * m_wavefrontSize * m_numWavefronts );
995  m_linkStrength.resize( m_maxBatchesWithinWave * m_wavefrontSize * m_numWavefronts );
996  m_linksMassLSC.resize( m_maxBatchesWithinWave * m_wavefrontSize * m_numWavefronts );
997  m_linksRestLengthSquared.resize( m_maxBatchesWithinWave * m_wavefrontSize * m_numWavefronts );
998  m_linksRestLength.resize( m_maxBatchesWithinWave * m_wavefrontSize * m_numWavefronts );
999  m_linksMaterialLinearStiffnessCoefficient.resize( m_maxBatchesWithinWave * m_wavefrontSize * m_numWavefronts );
1000 
1001  // Then re-order links into wavefront blocks
1002 
1003  // Total number of wavefronts moved. This will decide the ordering of sorted wavefronts.
1004  int wavefrontCount = 0;
1005 
1006  // Iterate over batches of wavefronts, then wavefronts in the batch
1007  for( int batchIndex = 0; batchIndex < numBatches; ++batchIndex )
1008  {
1009  btAlignedObjectArray <int> &batch( wavefrontBatches[batchIndex] );
1010  int wavefrontsInBatch = batch.size();
1011 
1012 
1013  for( int wavefrontIndex = 0; wavefrontIndex < wavefrontsInBatch; ++wavefrontIndex )
1014  {
1015 
1016  int originalWavefrontIndex = batch[wavefrontIndex];
1017  btAlignedObjectArray< int > &wavefrontVertices( verticesForWavefronts[originalWavefrontIndex] );
1018  int verticesUsedByWavefront = wavefrontVertices.size();
1019 
1020  // Copy the set of vertices into the correctly structured array for use on the device
1021  // Fill the non-vertices with -1s
1022  // so we can mask out those reads
1023  for( int vertex = 0; vertex < verticesUsedByWavefront; ++vertex )
1024  {
1025  m_wavefrontVerticesGlobalAddresses[m_maxVerticesWithinWave * wavefrontCount + vertex] = wavefrontVertices[vertex];
1026  }
1027  for( int vertex = verticesUsedByWavefront; vertex < m_maxVerticesWithinWave; ++vertex )
1028  {
1029  m_wavefrontVerticesGlobalAddresses[m_maxVerticesWithinWave * wavefrontCount + vertex] = -1;
1030  }
1031 
1032  // Obtain the set of batches within the current wavefront
1033  btAlignedObjectArray < btAlignedObjectArray <int> > &batchesWithinWavefront( batchesWithinWaves[originalWavefrontIndex] );
1034  // Set the size of the batches for use in the solver, correctly ordered
1035  NumBatchesVerticesPair batchesAndVertices;
1036  batchesAndVertices.numBatches = batchesWithinWavefront.size();
1037  batchesAndVertices.numVertices = verticesUsedByWavefront;
1038  m_numBatchesAndVerticesWithinWaves[wavefrontCount] = batchesAndVertices;
1039 
1040 
1041  // Now iterate over batches within the wavefront to structure the links correctly
1042  for( int wavefrontBatch = 0; wavefrontBatch < batchesWithinWavefront.size(); ++wavefrontBatch )
1043  {
1044  btAlignedObjectArray <int> &linksInBatch( batchesWithinWavefront[wavefrontBatch] );
1045  int wavefrontBatchSize = linksInBatch.size();
1046 
1047  int batchAddressInTarget = m_maxBatchesWithinWave * m_wavefrontSize * wavefrontCount + m_wavefrontSize * wavefrontBatch;
1048 
1049  for( int linkIndex = 0; linkIndex < wavefrontBatchSize; ++linkIndex )
1050  {
1051  int originalLinkAddress = linksInBatch[linkIndex];
1052  // Reorder simple arrays trivially
1053  m_links[batchAddressInTarget + linkIndex] = m_links_Backup[originalLinkAddress];
1054  m_linkStrength[batchAddressInTarget + linkIndex] = m_linkStrength_Backup[originalLinkAddress];
1055  m_linksMassLSC[batchAddressInTarget + linkIndex] = m_linksMassLSC_Backup[originalLinkAddress];
1056  m_linksRestLengthSquared[batchAddressInTarget + linkIndex] = m_linksRestLengthSquared_Backup[originalLinkAddress];
1057  m_linksRestLength[batchAddressInTarget + linkIndex] = m_linksRestLength_Backup[originalLinkAddress];
1058  m_linksMaterialLinearStiffnessCoefficient[batchAddressInTarget + linkIndex] = m_linksMaterialLinearStiffnessCoefficient_Backup[originalLinkAddress];
1059 
1060  // The local address is more complicated. We need to work out where a given vertex will end up
1061  // by searching the set of vertices for this link and using the index as the local address
1062  btSoftBodyLinkData::LinkNodePair localPair;
1063  btSoftBodyLinkData::LinkNodePair globalPair = m_links[batchAddressInTarget + linkIndex];
1064  localPair.vertex0 = wavefrontVertices.findLinearSearch( globalPair.vertex0 );
1065  localPair.vertex1 = wavefrontVertices.findLinearSearch( globalPair.vertex1 );
1066  m_linkVerticesLocalAddresses[batchAddressInTarget + linkIndex] = localPair;
1067  }
1068  for( int linkIndex = wavefrontBatchSize; linkIndex < m_wavefrontSize; ++linkIndex )
1069  {
1070  // Put 0s into these arrays for padding for cleanliness
1071  m_links[batchAddressInTarget + linkIndex] = btSoftBodyLinkData::LinkNodePair(0, 0);
1072  m_linkStrength[batchAddressInTarget + linkIndex] = 0.f;
1073  m_linksMassLSC[batchAddressInTarget + linkIndex] = 0.f;
1074  m_linksRestLengthSquared[batchAddressInTarget + linkIndex] = 0.f;
1075  m_linksRestLength[batchAddressInTarget + linkIndex] = 0.f;
1076  m_linksMaterialLinearStiffnessCoefficient[batchAddressInTarget + linkIndex] = 0.f;
1077 
1078 
1079  // For local addresses of junk data choose a set of addresses just above the range of valid ones
1080  // and cycling tyhrough % 16 so that we don't have bank conficts between all dud addresses
1081  // The valid addresses will do scatter and gather in the valid range, the junk ones should happily work
1082  // off the end of that range so we need no control
1083  btSoftBodyLinkData::LinkNodePair localPair;
1084  localPair.vertex0 = verticesUsedByWavefront + (linkIndex % 16);
1085  localPair.vertex1 = verticesUsedByWavefront + (linkIndex % 16);
1086  m_linkVerticesLocalAddresses[batchAddressInTarget + linkIndex] = localPair;
1087  }
1088 
1089  }
1090 
1091 
1092  wavefrontCount++;
1093  }
1094 
1095 
1096  }
1097 
1098 } // void btSoftBodyLinkDataDX11SIMDAware::generateBatches()
1099 
1100 
1101 
Generated on Thu Oct 3 2013 18:25:07 for Bullet Collision Detection & Physics Library by   doxygen 1.8.4