Bullet Collision Detection & Physics Library
btSoftBodySolver_OpenCL.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_OpenCL.h"
21 #include "BulletSoftBody/btSoftBodySolverVertexBuffer.h"
22 #include "BulletSoftBody/btSoftBody.h"
23 #include "BulletSoftBody/btSoftBodyInternals.h"
24 #include "BulletCollision/CollisionShapes/btCapsuleShape.h"
25 #include "BulletCollision/CollisionShapes/btSphereShape.h"
26 #include "LinearMath/btQuickprof.h"
27 #include <limits.h>
28 #include "BulletCollision/CollisionDispatch/btCollisionObjectWrapper.h"
29 
30 #define BT_SUPPRESS_OPENCL_ASSERTS
31 
32 #ifdef USE_MINICL
33  #include "MiniCL/cl.h"
34 #else //USE_MINICL
35  #ifdef __APPLE__
36  #include <OpenCL/OpenCL.h>
37  #else
38  #include <CL/cl.h>
39  #endif //__APPLE__
40 #endif//USE_MINICL
41 
42 #define BT_DEFAULT_WORKGROUPSIZE 64
43 
44 
45 #define RELEASE_CL_KERNEL(kernelName) {if( kernelName ){ clReleaseKernel( kernelName ); kernelName = 0; }}
46 
47 
48 //CL_VERSION_1_1 seems broken on NVidia SDK so just disable it
49 
51 #define MSTRINGIFY(A) #A
52 static const char* PrepareLinksCLString =
53 #include "OpenCLC10/PrepareLinks.cl"
54 static const char* UpdatePositionsFromVelocitiesCLString =
55 #include "OpenCLC10/UpdatePositionsFromVelocities.cl"
56 static const char* SolvePositionsCLString =
57 #include "OpenCLC10/SolvePositions.cl"
58 static const char* UpdateNodesCLString =
59 #include "OpenCLC10/UpdateNodes.cl"
60 static const char* UpdatePositionsCLString =
61 #include "OpenCLC10/UpdatePositions.cl"
62 static const char* UpdateConstantsCLString =
63 #include "OpenCLC10/UpdateConstants.cl"
64 static const char* IntegrateCLString =
65 #include "OpenCLC10/Integrate.cl"
66 static const char* ApplyForcesCLString =
67 #include "OpenCLC10/ApplyForces.cl"
68 static const char* UpdateFixedVertexPositionsCLString =
69 #include "OpenCLC10/UpdateFixedVertexPositions.cl"
70 static const char* UpdateNormalsCLString =
71 #include "OpenCLC10/UpdateNormals.cl"
72 static const char* VSolveLinksCLString =
73 #include "OpenCLC10/VSolveLinks.cl"
74 static const char* SolveCollisionsAndUpdateVelocitiesCLString =
75 #include "OpenCLC10/SolveCollisionsAndUpdateVelocities.cl"
76 
77 
78 btSoftBodyVertexDataOpenCL::btSoftBodyVertexDataOpenCL( cl_command_queue queue, cl_context ctx) :
79  m_queue(queue),
80  m_clClothIdentifier( queue, ctx, &m_clothIdentifier, false ),
81  m_clVertexPosition( queue, ctx, &m_vertexPosition, false ),
82  m_clVertexPreviousPosition( queue, ctx, &m_vertexPreviousPosition, false ),
83  m_clVertexVelocity( queue, ctx, &m_vertexVelocity, false ),
84  m_clVertexForceAccumulator( queue, ctx, &m_vertexForceAccumulator, false ),
85  m_clVertexNormal( queue, ctx, &m_vertexNormal, false ),
86  m_clVertexInverseMass( queue, ctx, &m_vertexInverseMass, false ),
87  m_clVertexArea( queue, ctx, &m_vertexArea, false ),
88  m_clVertexTriangleCount( queue, ctx, &m_vertexTriangleCount, false )
89 {
90 }
91 
92 btSoftBodyVertexDataOpenCL::~btSoftBodyVertexDataOpenCL()
93 {
94 
95 }
96 
97 bool btSoftBodyVertexDataOpenCL::onAccelerator()
98 {
99  return m_onGPU;
100 }
101 
102 bool btSoftBodyVertexDataOpenCL::moveToAccelerator()
103 {
104  bool success = true;
105  success = success && m_clClothIdentifier.moveToGPU();
106  success = success && m_clVertexPosition.moveToGPU();
107  success = success && m_clVertexPreviousPosition.moveToGPU();
108  success = success && m_clVertexVelocity.moveToGPU();
109  success = success && m_clVertexForceAccumulator.moveToGPU();
110  success = success && m_clVertexNormal.moveToGPU();
111  success = success && m_clVertexInverseMass.moveToGPU();
112  success = success && m_clVertexArea.moveToGPU();
113  success = success && m_clVertexTriangleCount.moveToGPU();
114 
115  if( success )
116  m_onGPU = true;
117 
118  return success;
119 }
120 
121 bool btSoftBodyVertexDataOpenCL::moveFromAccelerator(bool bCopy, bool bCopyMinimum)
122 {
123  bool success = true;
124 
125  if (!bCopy)
126  {
127  success = success && m_clClothIdentifier.moveFromGPU();
128  success = success && m_clVertexPosition.moveFromGPU();
129  success = success && m_clVertexPreviousPosition.moveFromGPU();
130  success = success && m_clVertexVelocity.moveFromGPU();
131  success = success && m_clVertexForceAccumulator.moveFromGPU();
132  success = success && m_clVertexNormal.moveFromGPU();
133  success = success && m_clVertexInverseMass.moveFromGPU();
134  success = success && m_clVertexArea.moveFromGPU();
135  success = success && m_clVertexTriangleCount.moveFromGPU();
136  }
137  else
138  {
139  if (bCopyMinimum)
140  {
141  success = success && m_clVertexPosition.copyFromGPU();
142  success = success && m_clVertexNormal.copyFromGPU();
143  }
144  else
145  {
146  success = success && m_clClothIdentifier.copyFromGPU();
147  success = success && m_clVertexPosition.copyFromGPU();
148  success = success && m_clVertexPreviousPosition.copyFromGPU();
149  success = success && m_clVertexVelocity.copyFromGPU();
150  success = success && m_clVertexForceAccumulator.copyFromGPU();
151  success = success && m_clVertexNormal.copyFromGPU();
152  success = success && m_clVertexInverseMass.copyFromGPU();
153  success = success && m_clVertexArea.copyFromGPU();
154  success = success && m_clVertexTriangleCount.copyFromGPU();
155  }
156  }
157 
158  if( success )
159  m_onGPU = true;
160 
161  return success;
162 }
163 
164 btSoftBodyLinkDataOpenCL::btSoftBodyLinkDataOpenCL(cl_command_queue queue, cl_context ctx)
165 :m_cqCommandQue(queue),
166  m_clLinks( queue, ctx, &m_links, false ),
167  m_clLinkStrength( queue, ctx, &m_linkStrength, false ),
168  m_clLinksMassLSC( queue, ctx, &m_linksMassLSC, false ),
169  m_clLinksRestLengthSquared( queue, ctx, &m_linksRestLengthSquared, false ),
170  m_clLinksCLength( queue, ctx, &m_linksCLength, false ),
171  m_clLinksLengthRatio( queue, ctx, &m_linksLengthRatio, false ),
172  m_clLinksRestLength( queue, ctx, &m_linksRestLength, false ),
173  m_clLinksMaterialLinearStiffnessCoefficient( queue, ctx, &m_linksMaterialLinearStiffnessCoefficient, false )
174 {
175 }
176 
177 btSoftBodyLinkDataOpenCL::~btSoftBodyLinkDataOpenCL()
178 {
179 }
180 
181 static Vectormath::Aos::Vector3 toVector3( const btVector3 &vec )
182 {
183  Vectormath::Aos::Vector3 outVec( vec.getX(), vec.getY(), vec.getZ() );
184  return outVec;
185 }
186 
188 void btSoftBodyLinkDataOpenCL::createLinks( int numLinks )
189 {
190  int previousSize = m_links.size();
191  int newSize = previousSize + numLinks;
192 
193  btSoftBodyLinkData::createLinks( numLinks );
194 
195  // Resize the link addresses array as well
196  m_linkAddresses.resize( newSize );
197 }
198 
200 void btSoftBodyLinkDataOpenCL::setLinkAt(
201  const LinkDescription &link,
202  int linkIndex )
203 {
204  btSoftBodyLinkData::setLinkAt( link, linkIndex );
205 
206  // Set the link index correctly for initialisation
207  m_linkAddresses[linkIndex] = linkIndex;
208 }
209 
210 bool btSoftBodyLinkDataOpenCL::onAccelerator()
211 {
212  return m_onGPU;
213 }
214 
215 bool btSoftBodyLinkDataOpenCL::moveToAccelerator()
216 {
217  bool success = true;
218  success = success && m_clLinks.moveToGPU();
219  success = success && m_clLinkStrength.moveToGPU();
220  success = success && m_clLinksMassLSC.moveToGPU();
221  success = success && m_clLinksRestLengthSquared.moveToGPU();
222  success = success && m_clLinksCLength.moveToGPU();
223  success = success && m_clLinksLengthRatio.moveToGPU();
224  success = success && m_clLinksRestLength.moveToGPU();
225  success = success && m_clLinksMaterialLinearStiffnessCoefficient.moveToGPU();
226 
227  if( success ) {
228  m_onGPU = true;
229  }
230 
231  return success;
232 }
233 
234 bool btSoftBodyLinkDataOpenCL::moveFromAccelerator()
235 {
236  bool success = true;
237  success = success && m_clLinks.moveFromGPU();
238  success = success && m_clLinkStrength.moveFromGPU();
239  success = success && m_clLinksMassLSC.moveFromGPU();
240  success = success && m_clLinksRestLengthSquared.moveFromGPU();
241  success = success && m_clLinksCLength.moveFromGPU();
242  success = success && m_clLinksLengthRatio.moveFromGPU();
243  success = success && m_clLinksRestLength.moveFromGPU();
244  success = success && m_clLinksMaterialLinearStiffnessCoefficient.moveFromGPU();
245 
246  if( success ) {
247  m_onGPU = false;
248  }
249 
250  return success;
251 }
252 
259 void btSoftBodyLinkDataOpenCL::generateBatches()
260 {
261  int numLinks = getNumLinks();
262 
263  // Do the graph colouring here temporarily
264  btAlignedObjectArray< int > batchValues;
265  batchValues.resize( numLinks, 0 );
266 
267  // Find the maximum vertex value internally for now
268  int maxVertex = 0;
269  for( int linkIndex = 0; linkIndex < numLinks; ++linkIndex )
270  {
271  int vertex0 = getVertexPair(linkIndex).vertex0;
272  int vertex1 = getVertexPair(linkIndex).vertex1;
273  if( vertex0 > maxVertex )
274  maxVertex = vertex0;
275  if( vertex1 > maxVertex )
276  maxVertex = vertex1;
277  }
278  int numVertices = maxVertex + 1;
279 
280  // Set of lists, one for each node, specifying which colours are connected
281  // to that node.
282  // No two edges into a node can share a colour.
283  btAlignedObjectArray< btAlignedObjectArray< int > > vertexConnectedColourLists;
284  vertexConnectedColourLists.resize(numVertices);
285 
286  // Simple algorithm that chooses the lowest batch number
287  // that none of the links attached to either of the connected
288  // nodes is in
289  for( int linkIndex = 0; linkIndex < numLinks; ++linkIndex )
290  {
291  int linkLocation = m_linkAddresses[linkIndex];
292 
293  int vertex0 = getVertexPair(linkLocation).vertex0;
294  int vertex1 = getVertexPair(linkLocation).vertex1;
295 
296  // Get the two node colour lists
297  btAlignedObjectArray< int > &colourListVertex0( vertexConnectedColourLists[vertex0] );
298  btAlignedObjectArray< int > &colourListVertex1( vertexConnectedColourLists[vertex1] );
299 
300  // Choose the minimum colour that is in neither list
301  int colour = 0;
302  while( colourListVertex0.findLinearSearch(colour) != colourListVertex0.size() || colourListVertex1.findLinearSearch(colour) != colourListVertex1.size() )
303  ++colour;
304  // i should now be the minimum colour in neither list
305  // Add to the two lists so that future edges don't share
306  // And store the colour against this edge
307 
308  colourListVertex0.push_back(colour);
309  colourListVertex1.push_back(colour);
310  batchValues[linkIndex] = colour;
311  }
312 
313  // Check the colour counts
314  btAlignedObjectArray< int > batchCounts;
315  for( int i = 0; i < numLinks; ++i )
316  {
317  int batch = batchValues[i];
318  if( batch >= batchCounts.size() )
319  batchCounts.push_back(1);
320  else
321  ++(batchCounts[batch]);
322  }
323 
324  m_batchStartLengths.resize(batchCounts.size());
325  if( m_batchStartLengths.size() > 0 )
326  {
327  m_batchStartLengths.resize(batchCounts.size());
328  m_batchStartLengths[0] = BatchPair(0, 0);
329 
330  int sum = 0;
331  for( int batchIndex = 0; batchIndex < batchCounts.size(); ++batchIndex )
332  {
333  m_batchStartLengths[batchIndex].start = sum;
334  m_batchStartLengths[batchIndex].length = batchCounts[batchIndex];
335  sum += batchCounts[batchIndex];
336  }
337  }
338 
340  // Sort data based on batches
341 
342  // Create source arrays by copying originals
343  btAlignedObjectArray<LinkNodePair> m_links_Backup(m_links);
344  btAlignedObjectArray<float> m_linkStrength_Backup(m_linkStrength);
345  btAlignedObjectArray<float> m_linksMassLSC_Backup(m_linksMassLSC);
346  btAlignedObjectArray<float> m_linksRestLengthSquared_Backup(m_linksRestLengthSquared);
347  btAlignedObjectArray<Vectormath::Aos::Vector3> m_linksCLength_Backup(m_linksCLength);
348  btAlignedObjectArray<float> m_linksLengthRatio_Backup(m_linksLengthRatio);
349  btAlignedObjectArray<float> m_linksRestLength_Backup(m_linksRestLength);
350  btAlignedObjectArray<float> m_linksMaterialLinearStiffnessCoefficient_Backup(m_linksMaterialLinearStiffnessCoefficient);
351 
352 
353  for( int batch = 0; batch < batchCounts.size(); ++batch )
354  batchCounts[batch] = 0;
355 
356  // Do sort as single pass into destination arrays
357  for( int linkIndex = 0; linkIndex < numLinks; ++linkIndex )
358  {
359  // To maintain locations run off the original link locations rather than the current position.
360  // It's not cache efficient, but as we run this rarely that should not matter.
361  // It's faster than searching the link location array for the current location and then updating it.
362  // The other alternative would be to unsort before resorting, but this is equivalent to doing that.
363  int linkLocation = m_linkAddresses[linkIndex];
364 
365  // Obtain batch and calculate target location for the
366  // next element in that batch, incrementing the batch counter
367  // afterwards
368  int batch = batchValues[linkIndex];
369  int newLocation = m_batchStartLengths[batch].start + batchCounts[batch];
370 
371  batchCounts[batch] = batchCounts[batch] + 1;
372  m_links[newLocation] = m_links_Backup[linkLocation];
373 #if 1
374  m_linkStrength[newLocation] = m_linkStrength_Backup[linkLocation];
375  m_linksMassLSC[newLocation] = m_linksMassLSC_Backup[linkLocation];
376  m_linksRestLengthSquared[newLocation] = m_linksRestLengthSquared_Backup[linkLocation];
377  m_linksLengthRatio[newLocation] = m_linksLengthRatio_Backup[linkLocation];
378  m_linksRestLength[newLocation] = m_linksRestLength_Backup[linkLocation];
379  m_linksMaterialLinearStiffnessCoefficient[newLocation] = m_linksMaterialLinearStiffnessCoefficient_Backup[linkLocation];
380 #endif
381  // Update the locations array to account for the moved entry
382  m_linkAddresses[linkIndex] = newLocation;
383  }
384 
385 
386 } // void generateBatches()
387 
388 
389 
390 
391 
392 btSoftBodyTriangleDataOpenCL::btSoftBodyTriangleDataOpenCL( cl_command_queue queue , cl_context ctx) :
393  m_queue( queue ),
394  m_clVertexIndices( queue, ctx, &m_vertexIndices, false ),
395  m_clArea( queue, ctx, &m_area, false ),
396  m_clNormal( queue, ctx, &m_normal, false )
397 {
398 }
399 
400 btSoftBodyTriangleDataOpenCL::~btSoftBodyTriangleDataOpenCL()
401 {
402 }
403 
405 void btSoftBodyTriangleDataOpenCL::createTriangles( int numTriangles )
406 {
407  int previousSize = getNumTriangles();
408  int newSize = previousSize + numTriangles;
409 
410  btSoftBodyTriangleData::createTriangles( numTriangles );
411 
412  // Resize the link addresses array as well
413  m_triangleAddresses.resize( newSize );
414 }
415 
417 void btSoftBodyTriangleDataOpenCL::setTriangleAt( const btSoftBodyTriangleData::TriangleDescription &triangle, int triangleIndex )
418 {
419  btSoftBodyTriangleData::setTriangleAt( triangle, triangleIndex );
420 
421  m_triangleAddresses[triangleIndex] = triangleIndex;
422 }
423 
424 bool btSoftBodyTriangleDataOpenCL::onAccelerator()
425 {
426  return m_onGPU;
427 }
428 
429 bool btSoftBodyTriangleDataOpenCL::moveToAccelerator()
430 {
431  bool success = true;
432  success = success && m_clVertexIndices.moveToGPU();
433  success = success && m_clArea.moveToGPU();
434  success = success && m_clNormal.moveToGPU();
435 
436  if( success )
437  m_onGPU = true;
438 
439  return success;
440 }
441 
442 bool btSoftBodyTriangleDataOpenCL::moveFromAccelerator()
443 {
444  bool success = true;
445  success = success && m_clVertexIndices.moveFromGPU();
446  success = success && m_clArea.moveFromGPU();
447  success = success && m_clNormal.moveFromGPU();
448 
449  if( success )
450  m_onGPU = true;
451 
452  return success;
453 }
454 
461 void btSoftBodyTriangleDataOpenCL::generateBatches()
462 {
463  int numTriangles = getNumTriangles();
464  if( numTriangles == 0 )
465  return;
466 
467  // Do the graph colouring here temporarily
468  btAlignedObjectArray< int > batchValues;
469  batchValues.resize( numTriangles );
470 
471  // Find the maximum vertex value internally for now
472  int maxVertex = 0;
473  for( int triangleIndex = 0; triangleIndex < numTriangles; ++triangleIndex )
474  {
475  int vertex0 = getVertexSet(triangleIndex).vertex0;
476  int vertex1 = getVertexSet(triangleIndex).vertex1;
477  int vertex2 = getVertexSet(triangleIndex).vertex2;
478 
479  if( vertex0 > maxVertex )
480  maxVertex = vertex0;
481  if( vertex1 > maxVertex )
482  maxVertex = vertex1;
483  if( vertex2 > maxVertex )
484  maxVertex = vertex2;
485  }
486  int numVertices = maxVertex + 1;
487 
488  // Set of lists, one for each node, specifying which colours are connected
489  // to that node.
490  // No two edges into a node can share a colour.
491  btAlignedObjectArray< btAlignedObjectArray< int > > vertexConnectedColourLists;
492  vertexConnectedColourLists.resize(numVertices);
493 
494 
495  //std::cout << "\n";
496  // Simple algorithm that chooses the lowest batch number
497  // that none of the faces attached to either of the connected
498  // nodes is in
499  for( int triangleIndex = 0; triangleIndex < numTriangles; ++triangleIndex )
500  {
501  // To maintain locations run off the original link locations rather than the current position.
502  // It's not cache efficient, but as we run this rarely that should not matter.
503  // It's faster than searching the link location array for the current location and then updating it.
504  // The other alternative would be to unsort before resorting, but this is equivalent to doing that.
505  int triangleLocation = m_triangleAddresses[triangleIndex];
506 
507  int vertex0 = getVertexSet(triangleLocation).vertex0;
508  int vertex1 = getVertexSet(triangleLocation).vertex1;
509  int vertex2 = getVertexSet(triangleLocation).vertex2;
510 
511  // Get the three node colour lists
512  btAlignedObjectArray< int > &colourListVertex0( vertexConnectedColourLists[vertex0] );
513  btAlignedObjectArray< int > &colourListVertex1( vertexConnectedColourLists[vertex1] );
514  btAlignedObjectArray< int > &colourListVertex2( vertexConnectedColourLists[vertex2] );
515 
516  // Choose the minimum colour that is in none of the lists
517  int colour = 0;
518  while(
519  colourListVertex0.findLinearSearch(colour) != colourListVertex0.size() ||
520  colourListVertex1.findLinearSearch(colour) != colourListVertex1.size() ||
521  colourListVertex2.findLinearSearch(colour) != colourListVertex2.size() )
522  {
523  ++colour;
524  }
525  // i should now be the minimum colour in neither list
526  // Add to the three lists so that future edges don't share
527  // And store the colour against this face
528  colourListVertex0.push_back(colour);
529  colourListVertex1.push_back(colour);
530  colourListVertex2.push_back(colour);
531 
532  batchValues[triangleIndex] = colour;
533  }
534 
535 
536  // Check the colour counts
537  btAlignedObjectArray< int > batchCounts;
538  for( int i = 0; i < numTriangles; ++i )
539  {
540  int batch = batchValues[i];
541  if( batch >= batchCounts.size() )
542  batchCounts.push_back(1);
543  else
544  ++(batchCounts[batch]);
545  }
546 
547 
548  m_batchStartLengths.resize(batchCounts.size());
549  m_batchStartLengths[0] = btSomePair(0,0);
550 
551 
552  int sum = 0;
553  for( int batchIndex = 0; batchIndex < batchCounts.size(); ++batchIndex )
554  {
555  m_batchStartLengths[batchIndex].first = sum;
556  m_batchStartLengths[batchIndex].second = batchCounts[batchIndex];
557  sum += batchCounts[batchIndex];
558  }
559 
561  // Sort data based on batches
562 
563  // Create source arrays by copying originals
564  btAlignedObjectArray<btSoftBodyTriangleData::TriangleNodeSet> m_vertexIndices_Backup(m_vertexIndices);
565  btAlignedObjectArray<float> m_area_Backup(m_area);
566  btAlignedObjectArray<Vectormath::Aos::Vector3> m_normal_Backup(m_normal);
567 
568 
569  for( int batch = 0; batch < batchCounts.size(); ++batch )
570  batchCounts[batch] = 0;
571 
572  // Do sort as single pass into destination arrays
573  for( int triangleIndex = 0; triangleIndex < numTriangles; ++triangleIndex )
574  {
575  // To maintain locations run off the original link locations rather than the current position.
576  // It's not cache efficient, but as we run this rarely that should not matter.
577  // It's faster than searching the link location array for the current location and then updating it.
578  // The other alternative would be to unsort before resorting, but this is equivalent to doing that.
579  int triangleLocation = m_triangleAddresses[triangleIndex];
580 
581  // Obtain batch and calculate target location for the
582  // next element in that batch, incrementing the batch counter
583  // afterwards
584  int batch = batchValues[triangleIndex];
585  int newLocation = m_batchStartLengths[batch].first + batchCounts[batch];
586 
587  batchCounts[batch] = batchCounts[batch] + 1;
588  m_vertexIndices[newLocation] = m_vertexIndices_Backup[triangleLocation];
589  m_area[newLocation] = m_area_Backup[triangleLocation];
590  m_normal[newLocation] = m_normal_Backup[triangleLocation];
591 
592  // Update the locations array to account for the moved entry
593  m_triangleAddresses[triangleIndex] = newLocation;
594  }
595 } // btSoftBodyTriangleDataOpenCL::generateBatches
596 
597 
598 
599 
600 
601 
602 
603 btOpenCLSoftBodySolver::btOpenCLSoftBodySolver(cl_command_queue queue, cl_context ctx, bool bUpdateAchchoredNodePos) :
604  m_linkData(queue, ctx),
605  m_vertexData(queue, ctx),
606  m_triangleData(queue, ctx),
607  m_defaultCLFunctions(queue, ctx),
608  m_currentCLFunctions(&m_defaultCLFunctions),
609  m_clPerClothAcceleration(queue, ctx, &m_perClothAcceleration, true ),
610  m_clPerClothWindVelocity(queue, ctx, &m_perClothWindVelocity, true ),
611  m_clPerClothDampingFactor(queue,ctx, &m_perClothDampingFactor, true ),
612  m_clPerClothVelocityCorrectionCoefficient(queue, ctx,&m_perClothVelocityCorrectionCoefficient, true ),
613  m_clPerClothLiftFactor(queue, ctx,&m_perClothLiftFactor, true ),
614  m_clPerClothDragFactor(queue, ctx,&m_perClothDragFactor, true ),
615  m_clPerClothMediumDensity(queue, ctx,&m_perClothMediumDensity, true ),
616  m_clPerClothCollisionObjects( queue, ctx, &m_perClothCollisionObjects, true ),
617  m_clCollisionObjectDetails( queue, ctx, &m_collisionObjectDetails, true ),
618  m_clPerClothFriction( queue, ctx, &m_perClothFriction, false ),
619  m_clAnchorPosition( queue, ctx, &m_anchorPosition, true ),
620  m_clAnchorIndex( queue, ctx, &m_anchorIndex, true),
621  m_cqCommandQue( queue ),
622  m_cxMainContext(ctx),
623  m_defaultWorkGroupSize(BT_DEFAULT_WORKGROUPSIZE),
624  m_bUpdateAnchoredNodePos(bUpdateAchchoredNodePos)
625 {
626 
627  // Initial we will clearly need to update solver constants
628  // For now this is global for the cloths linked with this solver - we should probably make this body specific
629  // for performance in future once we understand more clearly when constants need to be updated
630  m_updateSolverConstants = true;
631 
632  m_shadersInitialized = false;
633 
634  m_prepareLinksKernel = 0;
635  m_solvePositionsFromLinksKernel = 0;
636  m_updateConstantsKernel = 0;
637  m_integrateKernel = 0;
638  m_addVelocityKernel = 0;
639  m_updatePositionsFromVelocitiesKernel = 0;
640  m_updateVelocitiesFromPositionsWithoutVelocitiesKernel = 0;
641  m_updateVelocitiesFromPositionsWithVelocitiesKernel = 0;
642  m_vSolveLinksKernel = 0;
643  m_solveCollisionsAndUpdateVelocitiesKernel = 0;
644  m_resetNormalsAndAreasKernel = 0;
645  m_updateSoftBodiesKernel = 0;
646  m_normalizeNormalsAndAreasKernel = 0;
647  m_outputToVertexArrayKernel = 0;
648  m_applyForcesKernel = 0;
649  m_updateFixedVertexPositionsKernel = 0;
650 }
651 
652 btOpenCLSoftBodySolver::~btOpenCLSoftBodySolver()
653 {
654  releaseKernels();
655 }
656 
657 void btOpenCLSoftBodySolver::releaseKernels()
658 {
659  RELEASE_CL_KERNEL( m_prepareLinksKernel );
660  RELEASE_CL_KERNEL( m_solvePositionsFromLinksKernel );
661  RELEASE_CL_KERNEL( m_updateConstantsKernel );
662  RELEASE_CL_KERNEL( m_integrateKernel );
663  RELEASE_CL_KERNEL( m_addVelocityKernel );
664  RELEASE_CL_KERNEL( m_updatePositionsFromVelocitiesKernel );
665  RELEASE_CL_KERNEL( m_updateVelocitiesFromPositionsWithoutVelocitiesKernel );
666  RELEASE_CL_KERNEL( m_updateVelocitiesFromPositionsWithVelocitiesKernel );
667  RELEASE_CL_KERNEL( m_vSolveLinksKernel );
668  RELEASE_CL_KERNEL( m_solveCollisionsAndUpdateVelocitiesKernel );
669  RELEASE_CL_KERNEL( m_resetNormalsAndAreasKernel );
670  RELEASE_CL_KERNEL( m_normalizeNormalsAndAreasKernel );
671  RELEASE_CL_KERNEL( m_outputToVertexArrayKernel );
672  RELEASE_CL_KERNEL( m_applyForcesKernel );
673  RELEASE_CL_KERNEL( m_updateFixedVertexPositionsKernel );
674 
675  m_shadersInitialized = false;
676 }
677 
678 void btOpenCLSoftBodySolver::copyBackToSoftBodies(bool bMove)
679 {
680 
681  // Move the vertex data back to the host first
682  m_vertexData.moveFromAccelerator(!bMove);
683 
684  // Loop over soft bodies, copying all the vertex positions back for each body in turn
685  for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex )
686  {
687  btOpenCLAcceleratedSoftBodyInterface *softBodyInterface = m_softBodySet[ softBodyIndex ];
688  btSoftBody *softBody = softBodyInterface->getSoftBody();
689 
690  int firstVertex = softBodyInterface->getFirstVertex();
691  int numVertices = softBodyInterface->getNumVertices();
692 
693  // Copy vertices from solver back into the softbody
694  for( int vertex = 0; vertex < numVertices; ++vertex )
695  {
696  using Vectormath::Aos::Point3;
697  Point3 vertexPosition( m_vertexData.getVertexPositions()[firstVertex + vertex] );
698  Point3 normal(m_vertexData.getNormal(firstVertex + vertex));
699 
700  softBody->m_nodes[vertex].m_x.setX( vertexPosition.getX() );
701  softBody->m_nodes[vertex].m_x.setY( vertexPosition.getY() );
702  softBody->m_nodes[vertex].m_x.setZ( vertexPosition.getZ() );
703 
704  softBody->m_nodes[vertex].m_n.setX( normal.getX() );
705  softBody->m_nodes[vertex].m_n.setY( normal.getY() );
706  softBody->m_nodes[vertex].m_n.setZ( normal.getZ() );
707  }
708  }
709 } // btOpenCLSoftBodySolver::copyBackToSoftBodies
710 
711 void btOpenCLSoftBodySolver::optimize( btAlignedObjectArray< btSoftBody * > &softBodies, bool forceUpdate )
712 {
713  if( forceUpdate || m_softBodySet.size() != softBodies.size() )
714  {
715  // Have a change in the soft body set so update, reloading all the data
716  getVertexData().clear();
717  getTriangleData().clear();
718  getLinkData().clear();
719  m_softBodySet.resize(0);
720  m_anchorIndex.clear();
721 
722  int maxPiterations = 0;
723  int maxViterations = 0;
724 
725  for( int softBodyIndex = 0; softBodyIndex < softBodies.size(); ++softBodyIndex )
726  {
727  btSoftBody *softBody = softBodies[ softBodyIndex ];
728  using Vectormath::Aos::Matrix3;
729  using Vectormath::Aos::Point3;
730 
731  // Create SoftBody that will store the information within the solver
732  btOpenCLAcceleratedSoftBodyInterface *newSoftBody = new btOpenCLAcceleratedSoftBodyInterface( softBody );
733  m_softBodySet.push_back( newSoftBody );
734 
735  m_perClothAcceleration.push_back( toVector3(softBody->getWorldInfo()->m_gravity) );
736  m_perClothDampingFactor.push_back(softBody->m_cfg.kDP);
737  m_perClothVelocityCorrectionCoefficient.push_back( softBody->m_cfg.kVCF );
738  m_perClothLiftFactor.push_back( softBody->m_cfg.kLF );
739  m_perClothDragFactor.push_back( softBody->m_cfg.kDG );
740  m_perClothMediumDensity.push_back(softBody->getWorldInfo()->air_density);
741  // Simple init values. Actually we'll put 0 and -1 into them at the appropriate time
742  m_perClothFriction.push_back(softBody->m_cfg.kDF);
743  m_perClothCollisionObjects.push_back( CollisionObjectIndices(-1, -1) );
744 
745  // Add space for new vertices and triangles in the default solver for now
746  // TODO: Include space here for tearing too later
747  int firstVertex = getVertexData().getNumVertices();
748  int numVertices = softBody->m_nodes.size();
749  int maxVertices = numVertices;
750  // Allocate space for new vertices in all the vertex arrays
751  getVertexData().createVertices( maxVertices, softBodyIndex );
752 
753  int firstTriangle = getTriangleData().getNumTriangles();
754  int numTriangles = softBody->m_faces.size();
755  int maxTriangles = numTriangles;
756  getTriangleData().createTriangles( maxTriangles );
757 
758  // Copy vertices from softbody into the solver
759  for( int vertex = 0; vertex < numVertices; ++vertex )
760  {
761  Point3 multPoint(softBody->m_nodes[vertex].m_x.getX(), softBody->m_nodes[vertex].m_x.getY(), softBody->m_nodes[vertex].m_x.getZ());
762  btSoftBodyVertexData::VertexDescription desc;
763 
764  // TODO: Position in the softbody might be pre-transformed
765  // or we may need to adapt for the pose.
766  //desc.setPosition( cloth.getMeshTransform()*multPoint );
767  desc.setPosition( multPoint );
768 
769  float vertexInverseMass = softBody->m_nodes[vertex].m_im;
770  desc.setInverseMass(vertexInverseMass);
771  getVertexData().setVertexAt( desc, firstVertex + vertex );
772 
773  m_anchorIndex.push_back(-1);
774  }
775 
776  // Copy triangles similarly
777  // We're assuming here that vertex indices are based on the firstVertex rather than the entire scene
778  for( int triangle = 0; triangle < numTriangles; ++triangle )
779  {
780  // Note that large array storage is relative to the array not to the cloth
781  // So we need to add firstVertex to each value
782  int vertexIndex0 = (softBody->m_faces[triangle].m_n[0] - &(softBody->m_nodes[0]));
783  int vertexIndex1 = (softBody->m_faces[triangle].m_n[1] - &(softBody->m_nodes[0]));
784  int vertexIndex2 = (softBody->m_faces[triangle].m_n[2] - &(softBody->m_nodes[0]));
785  btSoftBodyTriangleData::TriangleDescription newTriangle(vertexIndex0 + firstVertex, vertexIndex1 + firstVertex, vertexIndex2 + firstVertex);
786  getTriangleData().setTriangleAt( newTriangle, firstTriangle + triangle );
787 
788  // Increase vertex triangle counts for this triangle
789  getVertexData().getTriangleCount(newTriangle.getVertexSet().vertex0)++;
790  getVertexData().getTriangleCount(newTriangle.getVertexSet().vertex1)++;
791  getVertexData().getTriangleCount(newTriangle.getVertexSet().vertex2)++;
792  }
793 
794  int firstLink = getLinkData().getNumLinks();
795  int numLinks = softBody->m_links.size();
796 // int maxLinks = numLinks;
797 
798  // Allocate space for the links
799  getLinkData().createLinks( numLinks );
800 
801  // Add the links
802  for( int link = 0; link < numLinks; ++link )
803  {
804  int vertexIndex0 = softBody->m_links[link].m_n[0] - &(softBody->m_nodes[0]);
805  int vertexIndex1 = softBody->m_links[link].m_n[1] - &(softBody->m_nodes[0]);
806 
807  btSoftBodyLinkData::LinkDescription newLink(vertexIndex0 + firstVertex, vertexIndex1 + firstVertex, softBody->m_links[link].m_material->m_kLST);
808  newLink.setLinkStrength(1.f);
809  getLinkData().setLinkAt(newLink, firstLink + link);
810  }
811 
812  newSoftBody->setFirstVertex( firstVertex );
813  newSoftBody->setFirstTriangle( firstTriangle );
814  newSoftBody->setNumVertices( numVertices );
815  newSoftBody->setMaxVertices( maxVertices );
816  newSoftBody->setNumTriangles( numTriangles );
817  newSoftBody->setMaxTriangles( maxTriangles );
818  newSoftBody->setFirstLink( firstLink );
819  newSoftBody->setNumLinks( numLinks );
820 
821  // Find maximum piterations and viterations
822  int piterations = softBody->m_cfg.piterations;
823 
824  if ( piterations > maxPiterations )
825  maxPiterations = piterations;
826 
827  int viterations = softBody->m_cfg.viterations;
828 
829  if ( viterations > maxViterations )
830  maxViterations = viterations;
831 
832  // zero mass
833  for( int vertex = 0; vertex < numVertices; ++vertex )
834  {
835  if ( softBody->m_nodes[vertex].m_im == 0 )
836  {
837  AnchorNodeInfoCL nodeInfo;
838  nodeInfo.clVertexIndex = firstVertex + vertex;
839  nodeInfo.pNode = &softBody->m_nodes[vertex];
840 
841  m_anchorNodeInfoArray.push_back(nodeInfo);
842  }
843  }
844 
845  // anchor position
846  if ( numVertices > 0 )
847  {
848  for ( int anchorIndex = 0; anchorIndex < softBody->m_anchors.size(); anchorIndex++ )
849  {
850  btSoftBody::Node* anchorNode = softBody->m_anchors[anchorIndex].m_node;
851  btSoftBody::Node* firstNode = &softBody->m_nodes[0];
852 
853  AnchorNodeInfoCL nodeInfo;
854  nodeInfo.clVertexIndex = firstVertex + (int)(anchorNode - firstNode);
855  nodeInfo.pNode = anchorNode;
856 
857  m_anchorNodeInfoArray.push_back(nodeInfo);
858  }
859  }
860  }
861 
862 
863  m_anchorPosition.clear();
864  m_anchorPosition.resize(m_anchorNodeInfoArray.size());
865 
866  for ( int anchorNode = 0; anchorNode < m_anchorNodeInfoArray.size(); anchorNode++ )
867  {
868  const AnchorNodeInfoCL& anchorNodeInfo = m_anchorNodeInfoArray[anchorNode];
869  m_anchorIndex[anchorNodeInfo.clVertexIndex] = anchorNode;
870  getVertexData().getInverseMass(anchorNodeInfo.clVertexIndex) = 0.0f;
871  }
872 
873  updateConstants(0.f);
874 
875  // set position and velocity iterations
876  setNumberOfPositionIterations(maxPiterations);
877  setNumberOfVelocityIterations(maxViterations);
878 
879  // set wind velocity
880  m_perClothWindVelocity.resize( m_softBodySet.size() );
881  for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex )
882  {
883  btSoftBody *softBody = m_softBodySet[softBodyIndex]->getSoftBody();
884  m_perClothWindVelocity[softBodyIndex] = toVector3(softBody->getWindVelocity());
885  }
886 
887  m_clPerClothWindVelocity.changedOnCPU();
888 
889  // generate batches
890  m_linkData.generateBatches();
891  m_triangleData.generateBatches();
892 
893  // Build the shaders to match the batching parameters
894  buildShaders();
895  }
896 }
897 
898 
899 btSoftBodyLinkData &btOpenCLSoftBodySolver::getLinkData()
900 {
901  // TODO: Consider setting link data to "changed" here
902  return m_linkData;
903 }
904 
905 btSoftBodyVertexData &btOpenCLSoftBodySolver::getVertexData()
906 {
907  // TODO: Consider setting vertex data to "changed" here
908  return m_vertexData;
909 }
910 
911 btSoftBodyTriangleData &btOpenCLSoftBodySolver::getTriangleData()
912 {
913  // TODO: Consider setting triangle data to "changed" here
914  return m_triangleData;
915 }
916 
917 void btOpenCLSoftBodySolver::resetNormalsAndAreas( int numVertices )
918 {
919  cl_int ciErrNum;
920  ciErrNum = clSetKernelArg(m_resetNormalsAndAreasKernel, 0, sizeof(numVertices), (void*)&numVertices); //oclCHECKERROR(ciErrNum, CL_SUCCESS);
921  ciErrNum = clSetKernelArg(m_resetNormalsAndAreasKernel, 1, sizeof(cl_mem), (void*)&m_vertexData.m_clVertexNormal.m_buffer);//oclCHECKERROR(ciErrNum, CL_SUCCESS);
922  ciErrNum = clSetKernelArg(m_resetNormalsAndAreasKernel, 2, sizeof(cl_mem), (void*)&m_vertexData.m_clVertexArea.m_buffer); //oclCHECKERROR(ciErrNum, CL_SUCCESS);
923  size_t numWorkItems = m_defaultWorkGroupSize*((numVertices + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize);
924 
925  if (numWorkItems)
926  {
927  ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue, m_resetNormalsAndAreasKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize, 0,0,0 );
928 
929  if( ciErrNum != CL_SUCCESS )
930  {
931  btAssert( 0 && "enqueueNDRangeKernel(m_resetNormalsAndAreasKernel)" );
932  }
933  }
934 
935 }
936 
937 void btOpenCLSoftBodySolver::normalizeNormalsAndAreas( int numVertices )
938 {
939  cl_int ciErrNum;
940 
941  ciErrNum = clSetKernelArg(m_normalizeNormalsAndAreasKernel, 0, sizeof(int),(void*) &numVertices);
942  ciErrNum = clSetKernelArg(m_normalizeNormalsAndAreasKernel, 1, sizeof(cl_mem), &m_vertexData.m_clVertexTriangleCount.m_buffer);
943  ciErrNum = clSetKernelArg(m_normalizeNormalsAndAreasKernel, 2, sizeof(cl_mem), &m_vertexData.m_clVertexNormal.m_buffer);
944  ciErrNum = clSetKernelArg(m_normalizeNormalsAndAreasKernel, 3, sizeof(cl_mem), &m_vertexData.m_clVertexArea.m_buffer);
945  size_t numWorkItems = m_defaultWorkGroupSize*((numVertices + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize);
946  if (numWorkItems)
947  {
948  ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue, m_normalizeNormalsAndAreasKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize, 0,0,0);
949  if( ciErrNum != CL_SUCCESS )
950  {
951  btAssert( 0 && "enqueueNDRangeKernel(m_normalizeNormalsAndAreasKernel)");
952  }
953  }
954 
955 }
956 
957 void btOpenCLSoftBodySolver::executeUpdateSoftBodies( int firstTriangle, int numTriangles )
958 {
959  cl_int ciErrNum;
960  ciErrNum = clSetKernelArg(m_updateSoftBodiesKernel, 0, sizeof(int), (void*) &firstTriangle);
961  ciErrNum = clSetKernelArg(m_updateSoftBodiesKernel, 1, sizeof(int), &numTriangles);
962  ciErrNum = clSetKernelArg(m_updateSoftBodiesKernel, 2, sizeof(cl_mem), &m_triangleData.m_clVertexIndices.m_buffer);
963  ciErrNum = clSetKernelArg(m_updateSoftBodiesKernel, 3, sizeof(cl_mem), &m_vertexData.m_clVertexPosition.m_buffer);
964  ciErrNum = clSetKernelArg(m_updateSoftBodiesKernel, 4, sizeof(cl_mem), &m_vertexData.m_clVertexNormal.m_buffer);
965  ciErrNum = clSetKernelArg(m_updateSoftBodiesKernel, 5, sizeof(cl_mem), &m_vertexData.m_clVertexArea.m_buffer);
966  ciErrNum = clSetKernelArg(m_updateSoftBodiesKernel, 6, sizeof(cl_mem), &m_triangleData.m_clNormal.m_buffer);
967  ciErrNum = clSetKernelArg(m_updateSoftBodiesKernel, 7, sizeof(cl_mem), &m_triangleData.m_clArea.m_buffer);
968 
969  size_t numWorkItems = m_defaultWorkGroupSize*((numTriangles + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize);
970  ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue, m_updateSoftBodiesKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize,0,0,0);
971  if( ciErrNum != CL_SUCCESS )
972  {
973  btAssert( 0 && "enqueueNDRangeKernel(m_normalizeNormalsAndAreasKernel)");
974  }
975 
976 }
977 
978 void btOpenCLSoftBodySolver::updateSoftBodies()
979 {
980  using namespace Vectormath::Aos;
981 
982 
983  int numVertices = m_vertexData.getNumVertices();
984 // int numTriangles = m_triangleData.getNumTriangles();
985 
986  // Ensure data is on accelerator
987  m_vertexData.moveToAccelerator();
988  m_triangleData.moveToAccelerator();
989 
990  resetNormalsAndAreas( numVertices );
991 
992 
993  // Go through triangle batches so updates occur correctly
994  for( int batchIndex = 0; batchIndex < m_triangleData.m_batchStartLengths.size(); ++batchIndex )
995  {
996 
997  int startTriangle = m_triangleData.m_batchStartLengths[batchIndex].first;
998  int numTriangles = m_triangleData.m_batchStartLengths[batchIndex].second;
999 
1000  executeUpdateSoftBodies( startTriangle, numTriangles );
1001  }
1002 
1003 
1004  normalizeNormalsAndAreas( numVertices );
1005 } // updateSoftBodies
1006 
1007 
1008 Vectormath::Aos::Vector3 btOpenCLSoftBodySolver::ProjectOnAxis( const Vectormath::Aos::Vector3 &v, const Vectormath::Aos::Vector3 &a )
1009 {
1010  return a*Vectormath::Aos::dot(v, a);
1011 }
1012 
1013 void btOpenCLSoftBodySolver::ApplyClampedForce( float solverdt, const Vectormath::Aos::Vector3 &force, const Vectormath::Aos::Vector3 &vertexVelocity, float inverseMass, Vectormath::Aos::Vector3 &vertexForce )
1014 {
1015  float dtInverseMass = solverdt*inverseMass;
1016  if( Vectormath::Aos::lengthSqr(force * dtInverseMass) > Vectormath::Aos::lengthSqr(vertexVelocity) )
1017  {
1018  vertexForce -= ProjectOnAxis( vertexVelocity, normalize( force ) )/dtInverseMass;
1019  } else {
1020  vertexForce += force;
1021  }
1022 }
1023 
1024 void btOpenCLSoftBodySolver::updateFixedVertexPositions()
1025 {
1026  // Ensure data is on accelerator
1027  m_vertexData.moveToAccelerator();
1028  m_clAnchorPosition.moveToGPU();
1029  m_clAnchorIndex.moveToGPU();
1030 
1031  cl_int ciErrNum ;
1032  int numVerts = m_vertexData.getNumVertices();
1033  ciErrNum = clSetKernelArg(m_updateFixedVertexPositionsKernel, 0, sizeof(int), &numVerts);
1034  ciErrNum = clSetKernelArg(m_updateFixedVertexPositionsKernel,1, sizeof(cl_mem), &m_clAnchorIndex.m_buffer);
1035  ciErrNum = clSetKernelArg(m_updateFixedVertexPositionsKernel,2, sizeof(cl_mem), &m_vertexData.m_clVertexPosition.m_buffer);
1036  ciErrNum = clSetKernelArg(m_updateFixedVertexPositionsKernel,3, sizeof(cl_mem), &m_clAnchorPosition.m_buffer);
1037 
1038  size_t numWorkItems = m_defaultWorkGroupSize*((m_vertexData.getNumVertices() + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize);
1039  if (numWorkItems)
1040  {
1041  ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,m_updateFixedVertexPositionsKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize, 0,0,0);
1042  if( ciErrNum != CL_SUCCESS )
1043  {
1044  btAssert( 0 && "enqueueNDRangeKernel(m_updateFixedVertexPositionsKernel)");
1045  }
1046  }
1047 
1048 }
1049 
1050 void btOpenCLSoftBodySolver::applyForces( float solverdt )
1051 {
1052  // Ensure data is on accelerator
1053  m_vertexData.moveToAccelerator();
1054  m_clPerClothAcceleration.moveToGPU();
1055  m_clPerClothLiftFactor.moveToGPU();
1056  m_clPerClothDragFactor.moveToGPU();
1057  m_clPerClothMediumDensity.moveToGPU();
1058  m_clPerClothWindVelocity.moveToGPU();
1059 
1060  cl_int ciErrNum ;
1061  int numVerts = m_vertexData.getNumVertices();
1062  ciErrNum = clSetKernelArg(m_applyForcesKernel, 0, sizeof(int), &numVerts);
1063  ciErrNum = clSetKernelArg(m_applyForcesKernel, 1, sizeof(float), &solverdt);
1064  float fl = FLT_EPSILON;
1065  ciErrNum = clSetKernelArg(m_applyForcesKernel, 2, sizeof(float), &fl);
1066  ciErrNum = clSetKernelArg(m_applyForcesKernel, 3, sizeof(cl_mem), &m_vertexData.m_clClothIdentifier.m_buffer);
1067  ciErrNum = clSetKernelArg(m_applyForcesKernel, 4, sizeof(cl_mem), &m_vertexData.m_clVertexNormal.m_buffer);
1068  ciErrNum = clSetKernelArg(m_applyForcesKernel, 5, sizeof(cl_mem), &m_vertexData.m_clVertexArea.m_buffer);
1069  ciErrNum = clSetKernelArg(m_applyForcesKernel, 6, sizeof(cl_mem), &m_vertexData.m_clVertexInverseMass.m_buffer);
1070  ciErrNum = clSetKernelArg(m_applyForcesKernel, 7, sizeof(cl_mem), &m_clPerClothLiftFactor.m_buffer);
1071  ciErrNum = clSetKernelArg(m_applyForcesKernel, 8 ,sizeof(cl_mem), &m_clPerClothDragFactor.m_buffer);
1072  ciErrNum = clSetKernelArg(m_applyForcesKernel, 9, sizeof(cl_mem), &m_clPerClothWindVelocity.m_buffer);
1073  ciErrNum = clSetKernelArg(m_applyForcesKernel,10, sizeof(cl_mem), &m_clPerClothAcceleration.m_buffer);
1074  ciErrNum = clSetKernelArg(m_applyForcesKernel,11, sizeof(cl_mem), &m_clPerClothMediumDensity.m_buffer);
1075  ciErrNum = clSetKernelArg(m_applyForcesKernel,12, sizeof(cl_mem), &m_vertexData.m_clVertexForceAccumulator.m_buffer);
1076  ciErrNum = clSetKernelArg(m_applyForcesKernel,13, sizeof(cl_mem), &m_vertexData.m_clVertexVelocity.m_buffer);
1077 
1078  size_t numWorkItems = m_defaultWorkGroupSize*((m_vertexData.getNumVertices() + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize);
1079  if (numWorkItems)
1080  {
1081  ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,m_applyForcesKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize, 0,0,0);
1082  if( ciErrNum != CL_SUCCESS )
1083  {
1084  btAssert( 0 && "enqueueNDRangeKernel(m_applyForcesKernel)");
1085  }
1086  }
1087 
1088 }
1089 
1093 void btOpenCLSoftBodySolver::integrate( float solverdt )
1094 {
1095  // Ensure data is on accelerator
1096  m_vertexData.moveToAccelerator();
1097 
1098  cl_int ciErrNum;
1099  int numVerts = m_vertexData.getNumVertices();
1100  ciErrNum = clSetKernelArg(m_integrateKernel, 0, sizeof(int), &numVerts);
1101  ciErrNum = clSetKernelArg(m_integrateKernel, 1, sizeof(float), &solverdt);
1102  ciErrNum = clSetKernelArg(m_integrateKernel, 2, sizeof(cl_mem), &m_vertexData.m_clVertexInverseMass.m_buffer);
1103  ciErrNum = clSetKernelArg(m_integrateKernel, 3, sizeof(cl_mem), &m_vertexData.m_clVertexPosition.m_buffer);
1104  ciErrNum = clSetKernelArg(m_integrateKernel, 4, sizeof(cl_mem), &m_vertexData.m_clVertexVelocity.m_buffer);
1105  ciErrNum = clSetKernelArg(m_integrateKernel, 5, sizeof(cl_mem), &m_vertexData.m_clVertexPreviousPosition.m_buffer);
1106  ciErrNum = clSetKernelArg(m_integrateKernel, 6, sizeof(cl_mem), &m_vertexData.m_clVertexForceAccumulator.m_buffer);
1107 
1108  size_t numWorkItems = m_defaultWorkGroupSize*((m_vertexData.getNumVertices() + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize);
1109  if (numWorkItems)
1110  {
1111  ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,m_integrateKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize,0,0,0);
1112  if( ciErrNum != CL_SUCCESS )
1113  {
1114  btAssert( 0 && "enqueueNDRangeKernel(m_integrateKernel)");
1115  }
1116  }
1117 
1118 }
1119 
1120 float btOpenCLSoftBodySolver::computeTriangleArea(
1121  const Vectormath::Aos::Point3 &vertex0,
1122  const Vectormath::Aos::Point3 &vertex1,
1123  const Vectormath::Aos::Point3 &vertex2 )
1124 {
1125  Vectormath::Aos::Vector3 a = vertex1 - vertex0;
1126  Vectormath::Aos::Vector3 b = vertex2 - vertex0;
1127  Vectormath::Aos::Vector3 crossProduct = cross(a, b);
1128  float area = length( crossProduct );
1129  return area;
1130 }
1131 
1132 
1133 void btOpenCLSoftBodySolver::updateBounds()
1134 {
1135  for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex )
1136  {
1137  btVector3 minBound(-1e30,-1e30,-1e30), maxBound(1e30,1e30,1e30);
1138  m_softBodySet[softBodyIndex]->updateBounds( minBound, maxBound );
1139  }
1140 
1141 } // btOpenCLSoftBodySolver::updateBounds
1142 
1143 
1144 void btOpenCLSoftBodySolver::updateConstants( float timeStep )
1145 {
1146 
1147  using namespace Vectormath::Aos;
1148 
1149  if( m_updateSolverConstants )
1150  {
1151  m_updateSolverConstants = false;
1152 
1153  // Will have to redo this if we change the structure (tear, maybe) or various other possible changes
1154 
1155  // Initialise link constants
1156  const int numLinks = m_linkData.getNumLinks();
1157  for( int linkIndex = 0; linkIndex < numLinks; ++linkIndex )
1158  {
1159  btSoftBodyLinkData::LinkNodePair &vertices( m_linkData.getVertexPair(linkIndex) );
1160  m_linkData.getRestLength(linkIndex) = length((m_vertexData.getPosition( vertices.vertex0 ) - m_vertexData.getPosition( vertices.vertex1 )));
1161  float invMass0 = m_vertexData.getInverseMass(vertices.vertex0);
1162  float invMass1 = m_vertexData.getInverseMass(vertices.vertex1);
1163  float linearStiffness = m_linkData.getLinearStiffnessCoefficient(linkIndex);
1164  float massLSC = (invMass0 + invMass1)/linearStiffness;
1165  m_linkData.getMassLSC(linkIndex) = massLSC;
1166  float restLength = m_linkData.getRestLength(linkIndex);
1167  float restLengthSquared = restLength*restLength;
1168  m_linkData.getRestLengthSquared(linkIndex) = restLengthSquared;
1169  }
1170  }
1171 
1172 }
1173 
1174 class QuickSortCompare
1175 {
1176  public:
1177 
1178  bool operator() ( const CollisionShapeDescription& a, const CollisionShapeDescription& b ) const
1179  {
1180  return ( a.softBodyIdentifier < b.softBodyIdentifier );
1181  }
1182 };
1183 
1184 
1188 void btOpenCLSoftBodySolver::prepareCollisionConstraints()
1189 {
1190  // First do a simple sort on the collision objects
1191  btAlignedObjectArray<int> numObjectsPerClothPrefixSum;
1192  btAlignedObjectArray<int> numObjectsPerCloth;
1193  numObjectsPerCloth.resize( m_softBodySet.size(), 0 );
1194  numObjectsPerClothPrefixSum.resize( m_softBodySet.size(), 0 );
1195 
1196 
1197 
1198  m_collisionObjectDetails.quickSort( QuickSortCompare() );
1199 
1200  if (!m_perClothCollisionObjects.size())
1201  return;
1202 
1203  // Generating indexing for perClothCollisionObjects
1204  // First clear the previous values with the "no collision object for cloth" constant
1205  for( int clothIndex = 0; clothIndex < m_perClothCollisionObjects.size(); ++clothIndex )
1206  {
1207  m_perClothCollisionObjects[clothIndex].firstObject = -1;
1208  m_perClothCollisionObjects[clothIndex].endObject = -1;
1209  }
1210  int currentCloth = 0;
1211  int startIndex = 0;
1212  for( int collisionObject = 0; collisionObject < m_collisionObjectDetails.size(); ++collisionObject )
1213  {
1214  int nextCloth = m_collisionObjectDetails[collisionObject].softBodyIdentifier;
1215  if( nextCloth != currentCloth )
1216  {
1217  // Changed cloth in the array
1218  // Set the end index and the range is what we need for currentCloth
1219  m_perClothCollisionObjects[currentCloth].firstObject = startIndex;
1220  m_perClothCollisionObjects[currentCloth].endObject = collisionObject;
1221  currentCloth = nextCloth;
1222  startIndex = collisionObject;
1223  }
1224  }
1225 
1226  // And update last cloth
1227  m_perClothCollisionObjects[currentCloth].firstObject = startIndex;
1228  m_perClothCollisionObjects[currentCloth].endObject = m_collisionObjectDetails.size();
1229 
1230 } // btOpenCLSoftBodySolver::prepareCollisionConstraints
1231 
1232 
1233 
1234 void btOpenCLSoftBodySolver::solveConstraints( float solverdt )
1235 {
1236 
1237  using Vectormath::Aos::Vector3;
1238  using Vectormath::Aos::Point3;
1239  using Vectormath::Aos::lengthSqr;
1240  using Vectormath::Aos::dot;
1241 
1242  // Prepare links
1243 // int numLinks = m_linkData.getNumLinks();
1244 // int numVertices = m_vertexData.getNumVertices();
1245 
1246  float kst = 1.f;
1247  float ti = 0.f;
1248 
1249 
1250  m_clPerClothDampingFactor.moveToGPU();
1251  m_clPerClothVelocityCorrectionCoefficient.moveToGPU();
1252 
1253 
1254  // Ensure data is on accelerator
1255  m_linkData.moveToAccelerator();
1256  m_vertexData.moveToAccelerator();
1257 
1258  prepareLinks();
1259 
1260 
1261 
1262  for( int iteration = 0; iteration < m_numberOfVelocityIterations ; ++iteration )
1263  {
1264  for( int i = 0; i < m_linkData.m_batchStartLengths.size(); ++i )
1265  {
1266  int startLink = m_linkData.m_batchStartLengths[i].start;
1267  int numLinks = m_linkData.m_batchStartLengths[i].length;
1268 
1269  solveLinksForVelocity( startLink, numLinks, kst );
1270  }
1271  }
1272 
1273 
1274  prepareCollisionConstraints();
1275 
1276  // Compute new positions from velocity
1277  // Also update the previous position so that our position computation is now based on the new position from the velocity solution
1278  // rather than based directly on the original positions
1279  if( m_numberOfVelocityIterations > 0 )
1280  {
1281  updateVelocitiesFromPositionsWithVelocities( 1.f/solverdt );
1282  } else {
1283  updateVelocitiesFromPositionsWithoutVelocities( 1.f/solverdt );
1284  }
1285 
1286  // Solve position
1287  for( int iteration = 0; iteration < m_numberOfPositionIterations ; ++iteration )
1288  {
1289  for( int i = 0; i < m_linkData.m_batchStartLengths.size(); ++i )
1290  {
1291  int startLink = m_linkData.m_batchStartLengths[i].start;
1292  int numLinks = m_linkData.m_batchStartLengths[i].length;
1293 
1294  solveLinksForPosition( startLink, numLinks, kst, ti );
1295  }
1296 
1297  } // for( int iteration = 0; iteration < m_numberOfPositionIterations ; ++iteration )
1298 
1299 
1300  // At this point assume that the force array is blank - we will overwrite it
1301  solveCollisionsAndUpdateVelocities( 1.f/solverdt );
1302 
1303 }
1304 
1305 
1307 // Kernel dispatches
1308 void btOpenCLSoftBodySolver::prepareLinks()
1309 {
1310  cl_int ciErrNum;
1311  int numLinks = m_linkData.getNumLinks();
1312  ciErrNum = clSetKernelArg(m_prepareLinksKernel,0, sizeof(int), &numLinks);
1313  ciErrNum = clSetKernelArg(m_prepareLinksKernel,1, sizeof(cl_mem), &m_linkData.m_clLinks.m_buffer);
1314  ciErrNum = clSetKernelArg(m_prepareLinksKernel,2, sizeof(cl_mem), &m_linkData.m_clLinksMassLSC.m_buffer);
1315  ciErrNum = clSetKernelArg(m_prepareLinksKernel,3, sizeof(cl_mem), &m_vertexData.m_clVertexPreviousPosition.m_buffer);
1316  ciErrNum = clSetKernelArg(m_prepareLinksKernel,4, sizeof(cl_mem), &m_linkData.m_clLinksLengthRatio.m_buffer);
1317  ciErrNum = clSetKernelArg(m_prepareLinksKernel,5, sizeof(cl_mem), &m_linkData.m_clLinksCLength.m_buffer);
1318 
1319  size_t numWorkItems = m_defaultWorkGroupSize*((m_linkData.getNumLinks() + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize);
1320  ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,m_prepareLinksKernel, 1 , NULL, &numWorkItems, &m_defaultWorkGroupSize,0,0,0);
1321  if( ciErrNum != CL_SUCCESS )
1322  {
1323  btAssert( 0 && "enqueueNDRangeKernel(m_prepareLinksKernel)");
1324  }
1325 
1326 }
1327 
1328 void btOpenCLSoftBodySolver::updatePositionsFromVelocities( float solverdt )
1329 {
1330  cl_int ciErrNum;
1331  int numVerts = m_vertexData.getNumVertices();
1332  ciErrNum = clSetKernelArg(m_updatePositionsFromVelocitiesKernel,0, sizeof(int), &numVerts);
1333  ciErrNum = clSetKernelArg(m_updatePositionsFromVelocitiesKernel,1, sizeof(float), &solverdt);
1334  ciErrNum = clSetKernelArg(m_updatePositionsFromVelocitiesKernel,2, sizeof(cl_mem), &m_vertexData.m_clVertexVelocity.m_buffer);
1335  ciErrNum = clSetKernelArg(m_updatePositionsFromVelocitiesKernel,3, sizeof(cl_mem), &m_vertexData.m_clVertexPreviousPosition.m_buffer);
1336  ciErrNum = clSetKernelArg(m_updatePositionsFromVelocitiesKernel,4, sizeof(cl_mem), &m_vertexData.m_clVertexPosition.m_buffer);
1337 
1338  size_t numWorkItems = m_defaultWorkGroupSize*((m_vertexData.getNumVertices() + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize);
1339  ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,m_updatePositionsFromVelocitiesKernel, 1, NULL, &numWorkItems,&m_defaultWorkGroupSize,0,0,0);
1340  if( ciErrNum != CL_SUCCESS )
1341  {
1342  btAssert( 0 && "enqueueNDRangeKernel(m_updatePositionsFromVelocitiesKernel)");
1343  }
1344 
1345 }
1346 
1347 void btOpenCLSoftBodySolver::solveLinksForPosition( int startLink, int numLinks, float kst, float ti )
1348 {
1349  cl_int ciErrNum;
1350  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,0, sizeof(int), &startLink);
1351  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,1, sizeof(int), &numLinks);
1352  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,2, sizeof(float), &kst);
1353  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,3, sizeof(float), &ti);
1354  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,4, sizeof(cl_mem), &m_linkData.m_clLinks.m_buffer);
1355  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,5, sizeof(cl_mem), &m_linkData.m_clLinksMassLSC.m_buffer);
1356  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,6, sizeof(cl_mem), &m_linkData.m_clLinksRestLengthSquared.m_buffer);
1357  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,7, sizeof(cl_mem), &m_vertexData.m_clVertexInverseMass.m_buffer);
1358  ciErrNum = clSetKernelArg(m_solvePositionsFromLinksKernel,8, sizeof(cl_mem), &m_vertexData.m_clVertexPosition.m_buffer);
1359 
1360  size_t numWorkItems = m_defaultWorkGroupSize*((numLinks + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize);
1361  ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,m_solvePositionsFromLinksKernel,1,NULL,&numWorkItems,&m_defaultWorkGroupSize,0,0,0);
1362  if( ciErrNum!= CL_SUCCESS )
1363  {
1364  btAssert( 0 && "enqueueNDRangeKernel(m_solvePositionsFromLinksKernel)");
1365  }
1366 
1367 } // solveLinksForPosition
1368 
1369 
1370 void btOpenCLSoftBodySolver::solveLinksForVelocity( int startLink, int numLinks, float kst )
1371 {
1372  cl_int ciErrNum;
1373  ciErrNum = clSetKernelArg(m_vSolveLinksKernel, 0, sizeof(int), &startLink);
1374  ciErrNum = clSetKernelArg(m_vSolveLinksKernel, 1, sizeof(int), &numLinks);
1375  ciErrNum = clSetKernelArg(m_vSolveLinksKernel, 2, sizeof(float), &kst);
1376  ciErrNum = clSetKernelArg(m_vSolveLinksKernel, 3, sizeof(cl_mem), &m_linkData.m_clLinks.m_buffer);
1377  ciErrNum = clSetKernelArg(m_vSolveLinksKernel, 4, sizeof(cl_mem), &m_linkData.m_clLinksLengthRatio.m_buffer);
1378  ciErrNum = clSetKernelArg(m_vSolveLinksKernel, 5, sizeof(cl_mem), &m_linkData.m_clLinksCLength.m_buffer);
1379  ciErrNum = clSetKernelArg(m_vSolveLinksKernel, 6, sizeof(cl_mem), &m_vertexData.m_clVertexInverseMass.m_buffer);
1380  ciErrNum = clSetKernelArg(m_vSolveLinksKernel, 7, sizeof(cl_mem), &m_vertexData.m_clVertexVelocity.m_buffer);
1381 
1382  size_t numWorkItems = m_defaultWorkGroupSize*((numLinks + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize);
1383  ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,m_vSolveLinksKernel,1,NULL,&numWorkItems, &m_defaultWorkGroupSize,0,0,0);
1384  if( ciErrNum != CL_SUCCESS )
1385  {
1386  btAssert( 0 && "enqueueNDRangeKernel(m_vSolveLinksKernel)");
1387  }
1388 
1389 }
1390 
1391 void btOpenCLSoftBodySolver::updateVelocitiesFromPositionsWithVelocities( float isolverdt )
1392 {
1393  cl_int ciErrNum;
1394  int numVerts = m_vertexData.getNumVertices();
1395  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithVelocitiesKernel,0, sizeof(int), &numVerts);
1396  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithVelocitiesKernel, 1, sizeof(float), &isolverdt);
1397  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithVelocitiesKernel, 2, sizeof(cl_mem), &m_vertexData.m_clVertexPosition.m_buffer);
1398  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithVelocitiesKernel, 3, sizeof(cl_mem), &m_vertexData.m_clVertexPreviousPosition.m_buffer);
1399  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithVelocitiesKernel, 4, sizeof(cl_mem), &m_vertexData.m_clClothIdentifier.m_buffer);
1400  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithVelocitiesKernel, 5, sizeof(cl_mem), &m_clPerClothVelocityCorrectionCoefficient.m_buffer);
1401  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithVelocitiesKernel, 6, sizeof(cl_mem), &m_clPerClothDampingFactor.m_buffer);
1402  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithVelocitiesKernel, 7, sizeof(cl_mem), &m_vertexData.m_clVertexVelocity.m_buffer);
1403  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithVelocitiesKernel, 8, sizeof(cl_mem), &m_vertexData.m_clVertexForceAccumulator.m_buffer);
1404 
1405  size_t numWorkItems = m_defaultWorkGroupSize*((m_vertexData.getNumVertices() + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize);
1406  ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,m_updateVelocitiesFromPositionsWithVelocitiesKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize,0,0,0);
1407  if( ciErrNum != CL_SUCCESS )
1408  {
1409  btAssert( 0 && "enqueueNDRangeKernel(m_updateVelocitiesFromPositionsWithVelocitiesKernel)");
1410  }
1411 
1412 
1413 } // updateVelocitiesFromPositionsWithVelocities
1414 
1415 void btOpenCLSoftBodySolver::updateVelocitiesFromPositionsWithoutVelocities( float isolverdt )
1416 {
1417  cl_int ciErrNum;
1418  int numVerts = m_vertexData.getNumVertices();
1419  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithoutVelocitiesKernel, 0, sizeof(int), &numVerts);
1420  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithoutVelocitiesKernel, 1, sizeof(float), &isolverdt);
1421  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithoutVelocitiesKernel, 2, sizeof(cl_mem),&m_vertexData.m_clVertexPosition.m_buffer);
1422  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithoutVelocitiesKernel, 3, sizeof(cl_mem),&m_vertexData.m_clVertexPreviousPosition.m_buffer);
1423  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithoutVelocitiesKernel, 4, sizeof(cl_mem),&m_vertexData.m_clClothIdentifier.m_buffer);
1424  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithoutVelocitiesKernel, 5, sizeof(cl_mem),&m_clPerClothDampingFactor.m_buffer);
1425  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithoutVelocitiesKernel, 6, sizeof(cl_mem),&m_vertexData.m_clVertexVelocity.m_buffer);
1426  ciErrNum = clSetKernelArg(m_updateVelocitiesFromPositionsWithoutVelocitiesKernel, 7, sizeof(cl_mem),&m_vertexData.m_clVertexForceAccumulator.m_buffer);
1427 
1428  size_t numWorkItems = m_defaultWorkGroupSize*((m_vertexData.getNumVertices() + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize);
1429  ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,m_updateVelocitiesFromPositionsWithoutVelocitiesKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize,0,0,0);
1430  if( ciErrNum != CL_SUCCESS )
1431  {
1432  btAssert( 0 && "enqueueNDRangeKernel(m_updateVelocitiesFromPositionsWithoutVelocitiesKernel)");
1433  }
1434 
1435 } // updateVelocitiesFromPositionsWithoutVelocities
1436 
1437 
1438 
1439 void btOpenCLSoftBodySolver::solveCollisionsAndUpdateVelocities( float isolverdt )
1440 {
1441  // Copy kernel parameters to GPU
1442  m_vertexData.moveToAccelerator();
1443  m_clPerClothFriction.moveToGPU();
1444  m_clPerClothDampingFactor.moveToGPU();
1445  m_clPerClothCollisionObjects.moveToGPU();
1446  m_clCollisionObjectDetails.moveToGPU();
1447 
1448  cl_int ciErrNum;
1449  int numVerts = m_vertexData.getNumVertices();
1450  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 0, sizeof(int), &numVerts);
1451  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 1, sizeof(int), &isolverdt);
1452  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 2, sizeof(cl_mem),&m_vertexData.m_clClothIdentifier.m_buffer);
1453  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 3, sizeof(cl_mem),&m_vertexData.m_clVertexPreviousPosition.m_buffer);
1454  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 4, sizeof(cl_mem),&m_clPerClothFriction.m_buffer);
1455  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 5, sizeof(cl_mem),&m_clPerClothDampingFactor.m_buffer);
1456  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 6, sizeof(cl_mem),&m_clPerClothCollisionObjects.m_buffer);
1457  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 7, sizeof(cl_mem),&m_clCollisionObjectDetails.m_buffer);
1458  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 8, sizeof(cl_mem),&m_vertexData.m_clVertexForceAccumulator.m_buffer);
1459  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 9, sizeof(cl_mem),&m_vertexData.m_clVertexVelocity.m_buffer);
1460  ciErrNum = clSetKernelArg(m_solveCollisionsAndUpdateVelocitiesKernel, 10, sizeof(cl_mem),&m_vertexData.m_clVertexPosition.m_buffer);
1461 
1462  size_t numWorkItems = m_defaultWorkGroupSize*((m_vertexData.getNumVertices() + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize);
1463  if (numWorkItems)
1464  {
1465  ciErrNum = clEnqueueNDRangeKernel(m_cqCommandQue,m_solveCollisionsAndUpdateVelocitiesKernel, 1, NULL, &numWorkItems, &m_defaultWorkGroupSize,0,0,0);
1466  if( ciErrNum != CL_SUCCESS )
1467  {
1468  btAssert( 0 && "enqueueNDRangeKernel(m_updateVelocitiesFromPositionsWithoutVelocitiesKernel)");
1469  }
1470  }
1471 
1472 } // btOpenCLSoftBodySolver::updateVelocitiesFromPositionsWithoutVelocities
1473 
1474 
1475 
1476 // End kernel dispatches
1478 
1479 
1480 void btSoftBodySolverOutputCLtoCPU::copySoftBodyToVertexBuffer( const btSoftBody * const softBody, btVertexBufferDescriptor *vertexBuffer )
1481 {
1482 
1483  btSoftBodySolver *solver = softBody->getSoftBodySolver();
1484  btAssert( solver->getSolverType() == btSoftBodySolver::CL_SOLVER || solver->getSolverType() == btSoftBodySolver::CL_SIMD_SOLVER );
1485  btOpenCLSoftBodySolver *dxSolver = static_cast< btOpenCLSoftBodySolver * >( solver );
1486 
1487  btOpenCLAcceleratedSoftBodyInterface* currentCloth = dxSolver->findSoftBodyInterface( softBody );
1488  btSoftBodyVertexDataOpenCL &vertexData( dxSolver->m_vertexData );
1489 
1490 
1491  const int firstVertex = currentCloth->getFirstVertex();
1492  const int lastVertex = firstVertex + currentCloth->getNumVertices();
1493 
1494  if( vertexBuffer->getBufferType() == btVertexBufferDescriptor::CPU_BUFFER )
1495  {
1496  const btCPUVertexBufferDescriptor *cpuVertexBuffer = static_cast< btCPUVertexBufferDescriptor* >(vertexBuffer);
1497  float *basePointer = cpuVertexBuffer->getBasePointer();
1498 
1499  vertexData.m_clVertexPosition.copyFromGPU();
1500  vertexData.m_clVertexNormal.copyFromGPU();
1501 
1502  if( vertexBuffer->hasVertexPositions() )
1503  {
1504  const int vertexOffset = cpuVertexBuffer->getVertexOffset();
1505  const int vertexStride = cpuVertexBuffer->getVertexStride();
1506  float *vertexPointer = basePointer + vertexOffset;
1507 
1508  for( int vertexIndex = firstVertex; vertexIndex < lastVertex; ++vertexIndex )
1509  {
1510  Vectormath::Aos::Point3 position = vertexData.getPosition(vertexIndex);
1511  *(vertexPointer + 0) = position.getX();
1512  *(vertexPointer + 1) = position.getY();
1513  *(vertexPointer + 2) = position.getZ();
1514  vertexPointer += vertexStride;
1515  }
1516  }
1517  if( vertexBuffer->hasNormals() )
1518  {
1519  const int normalOffset = cpuVertexBuffer->getNormalOffset();
1520  const int normalStride = cpuVertexBuffer->getNormalStride();
1521  float *normalPointer = basePointer + normalOffset;
1522 
1523  for( int vertexIndex = firstVertex; vertexIndex < lastVertex; ++vertexIndex )
1524  {
1525  Vectormath::Aos::Vector3 normal = vertexData.getNormal(vertexIndex);
1526  *(normalPointer + 0) = normal.getX();
1527  *(normalPointer + 1) = normal.getY();
1528  *(normalPointer + 2) = normal.getZ();
1529  normalPointer += normalStride;
1530  }
1531  }
1532  }
1533 
1534 } // btSoftBodySolverOutputCLtoCPU::outputToVertexBuffers
1535 
1536 
1537 
1538 cl_kernel CLFunctions::compileCLKernelFromString( const char* kernelSource, const char* kernelName, const char* additionalMacros ,const char* orgSrcFileNameForCaching)
1539 {
1540  printf("compiling kernelName: %s ",kernelName);
1541  cl_kernel kernel=0;
1542  cl_int ciErrNum;
1543  size_t program_length = strlen(kernelSource);
1544 
1545  cl_program m_cpProgram = clCreateProgramWithSource(m_cxMainContext, 1, (const char**)&kernelSource, &program_length, &ciErrNum);
1546 // oclCHECKERROR(ciErrNum, CL_SUCCESS);
1547 
1548  // Build the program with 'mad' Optimization option
1549 
1550 
1551 #ifdef MAC
1552  char* flags = "-cl-mad-enable -DMAC -DGUID_ARG";
1553 #else
1554  //const char* flags = "-DGUID_ARG= -fno-alias";
1555  const char* flags = "-DGUID_ARG= ";
1556 #endif
1557 
1558  char* compileFlags = new char[strlen(additionalMacros) + strlen(flags) + 5];
1559  sprintf(compileFlags, "%s %s", flags, additionalMacros);
1560  ciErrNum = clBuildProgram(m_cpProgram, 0, NULL, compileFlags, NULL, NULL);
1561  if (ciErrNum != CL_SUCCESS)
1562  {
1563  size_t numDevices;
1564  clGetProgramInfo( m_cpProgram, CL_PROGRAM_DEVICES, 0, 0, &numDevices );
1565  cl_device_id *devices = new cl_device_id[numDevices];
1566  clGetProgramInfo( m_cpProgram, CL_PROGRAM_DEVICES, numDevices, devices, &numDevices );
1567  for( int i = 0; i < 2; ++i )
1568  {
1569  char *build_log;
1570  size_t ret_val_size;
1571  clGetProgramBuildInfo(m_cpProgram, devices[i], CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
1572  build_log = new char[ret_val_size+1];
1573  clGetProgramBuildInfo(m_cpProgram, devices[i], CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
1574 
1575  // to be carefully, terminate with \0
1576  // there's no information in the reference whether the string is 0 terminated or not
1577  build_log[ret_val_size] = '\0';
1578 
1579 
1580  printf("Error in clBuildProgram, Line %u in file %s, Log: \n%s\n !!!\n\n", __LINE__, __FILE__, build_log);
1581  delete[] build_log;
1582  }
1583 #ifndef BT_SUPPRESS_OPENCL_ASSERTS
1584  btAssert(0);
1585 #endif //BT_SUPPRESS_OPENCL_ASSERTS
1586  m_kernelCompilationFailures++;
1587  return 0;
1588  }
1589 
1590 
1591  // Create the kernel
1592  kernel = clCreateKernel(m_cpProgram, kernelName, &ciErrNum);
1593  if (ciErrNum != CL_SUCCESS)
1594  {
1595  const char* msg = "";
1596  switch(ciErrNum)
1597  {
1598  case CL_INVALID_PROGRAM:
1599  msg = "Program is not a valid program object.";
1600  break;
1601  case CL_INVALID_PROGRAM_EXECUTABLE:
1602  msg = "There is no successfully built executable for program.";
1603  break;
1604  case CL_INVALID_KERNEL_NAME:
1605  msg = "kernel_name is not found in program.";
1606  break;
1607  case CL_INVALID_KERNEL_DEFINITION:
1608  msg = "the function definition for __kernel function given by kernel_name such as the number of arguments, the argument types are not the same for all devices for which the program executable has been built.";
1609  break;
1610  case CL_INVALID_VALUE:
1611  msg = "kernel_name is NULL.";
1612  break;
1613  case CL_OUT_OF_HOST_MEMORY:
1614  msg = "Failure to allocate resources required by the OpenCL implementation on the host.";
1615  break;
1616  default:
1617  {
1618  }
1619  }
1620 
1621  printf("Error in clCreateKernel for kernel '%s', error is \"%s\", Line %u in file %s !!!\n\n", kernelName, msg, __LINE__, __FILE__);
1622 
1623 #ifndef BT_SUPPRESS_OPENCL_ASSERTS
1624  btAssert(0);
1625 #endif //BT_SUPPRESS_OPENCL_ASSERTS
1626  m_kernelCompilationFailures++;
1627  return 0;
1628  }
1629 
1630  printf("ready. \n");
1631  delete [] compileFlags;
1632  if (!kernel)
1633  m_kernelCompilationFailures++;
1634  return kernel;
1635 
1636 }
1637 
1638 void btOpenCLSoftBodySolver::predictMotion( float timeStep )
1639 {
1640  // Clear the collision shape array for the next frame
1641  // Ensure that the DX11 ones are moved off the device so they will be updated correctly
1642  m_clCollisionObjectDetails.changedOnCPU();
1643  m_clPerClothCollisionObjects.changedOnCPU();
1644  m_collisionObjectDetails.clear();
1645 
1646  if ( m_bUpdateAnchoredNodePos )
1647  {
1648  // In OpenCL cloth solver, if softbody node has zero inverse mass(infinite mass) or anchor attached,
1649  // we need to update the node position in case the node or anchor is animated externally.
1650  // If there is no such node, we can eliminate the unnecessary CPU-to-GPU data trasferring.
1651  for ( int i = 0; i < m_anchorNodeInfoArray.size(); i++ )
1652  {
1653  const AnchorNodeInfoCL& anchorNodeInfo = m_anchorNodeInfoArray[i];
1654  btSoftBody::Node* node = anchorNodeInfo.pNode;
1655 
1656  using Vectormath::Aos::Point3;
1657  Point3 pos((float)node->m_x.getX(), (float)node->m_x.getY(), (float)node->m_x.getZ());
1658  m_anchorPosition[i] = pos;
1659  }
1660 
1661  if ( m_anchorNodeInfoArray.size() > 0 )
1662  m_clAnchorPosition.changedOnCPU();
1663 
1664  updateFixedVertexPositions();
1665  }
1666 
1667  {
1668  BT_PROFILE("applyForces");
1669  // Apply forces that we know about to the cloths
1670  applyForces( timeStep * getTimeScale() );
1671  }
1672 
1673  {
1674  BT_PROFILE("integrate");
1675  // Itegrate motion for all soft bodies dealt with by the solver
1676  integrate( timeStep * getTimeScale() );
1677  }
1678 
1679  {
1680  BT_PROFILE("updateBounds");
1681  updateBounds();
1682  }
1683  // End prediction work for solvers
1684 }
1685 
1686 static Vectormath::Aos::Transform3 toTransform3( const btTransform &transform )
1687 {
1688  Vectormath::Aos::Transform3 outTransform;
1689  outTransform.setCol(0, toVector3(transform.getBasis().getColumn(0)));
1690  outTransform.setCol(1, toVector3(transform.getBasis().getColumn(1)));
1691  outTransform.setCol(2, toVector3(transform.getBasis().getColumn(2)));
1692  outTransform.setCol(3, toVector3(transform.getOrigin()));
1693  return outTransform;
1694 }
1695 
1696 void btOpenCLAcceleratedSoftBodyInterface::updateBounds( const btVector3 &lowerBound, const btVector3 &upperBound )
1697 {
1698  float scalarMargin = (float)getSoftBody()->getCollisionShape()->getMargin();
1699  btVector3 vectorMargin( scalarMargin, scalarMargin, scalarMargin );
1700  m_softBody->m_bounds[0] = lowerBound - vectorMargin;
1701  m_softBody->m_bounds[1] = upperBound + vectorMargin;
1702 } // btOpenCLSoftBodySolver::btDX11AcceleratedSoftBodyInterface::updateBounds
1703 
1704 void btOpenCLSoftBodySolver::processCollision( btSoftBody*, btSoftBody* )
1705 {
1706 
1707 }
1708 
1709 // Add the collision object to the set to deal with for a particular soft body
1710 void btOpenCLSoftBodySolver::processCollision( btSoftBody *softBody, const btCollisionObjectWrapper* collisionObject )
1711 {
1712  int softBodyIndex = findSoftBodyIndex( softBody );
1713 
1714  if( softBodyIndex >= 0 )
1715  {
1716  const btCollisionShape *collisionShape = collisionObject->getCollisionShape();
1717  float friction = collisionObject->getCollisionObject()->getFriction();
1718  int shapeType = collisionShape->getShapeType();
1719  if( shapeType == CAPSULE_SHAPE_PROXYTYPE )
1720  {
1721  // Add to the list of expected collision objects
1722  CollisionShapeDescription newCollisionShapeDescription;
1723  newCollisionShapeDescription.softBodyIdentifier = softBodyIndex;
1724  newCollisionShapeDescription.collisionShapeType = shapeType;
1725  // TODO: May need to transpose this matrix either here or in HLSL
1726  newCollisionShapeDescription.shapeTransform = toTransform3(collisionObject->getWorldTransform());
1727  const btCapsuleShape *capsule = static_cast<const btCapsuleShape*>( collisionShape );
1728  newCollisionShapeDescription.radius = capsule->getRadius();
1729  newCollisionShapeDescription.halfHeight = capsule->getHalfHeight();
1730  newCollisionShapeDescription.margin = capsule->getMargin();
1731  newCollisionShapeDescription.upAxis = capsule->getUpAxis();
1732  newCollisionShapeDescription.friction = friction;
1733  const btRigidBody* body = static_cast< const btRigidBody* >( collisionObject->getCollisionObject() );
1734  newCollisionShapeDescription.linearVelocity = toVector3(body->getLinearVelocity());
1735  newCollisionShapeDescription.angularVelocity = toVector3(body->getAngularVelocity());
1736  m_collisionObjectDetails.push_back( newCollisionShapeDescription );
1737 
1738  }
1739  else {
1740 #ifdef _DEBUG
1741  printf("Unsupported collision shape type\n");
1742 #endif
1743  //btAssert(0 && "Unsupported collision shape type\n");
1744  }
1745  } else {
1746  btAssert(0 && "Unknown soft body");
1747  }
1748 } // btOpenCLSoftBodySolver::processCollision
1749 
1750 
1751 
1752 
1753 
1754 btOpenCLAcceleratedSoftBodyInterface* btOpenCLSoftBodySolver::findSoftBodyInterface( const btSoftBody* const softBody )
1755 {
1756  for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex )
1757  {
1758  btOpenCLAcceleratedSoftBodyInterface* softBodyInterface = m_softBodySet[softBodyIndex];
1759  if( softBodyInterface->getSoftBody() == softBody )
1760  return softBodyInterface;
1761  }
1762  return 0;
1763 }
1764 
1765 
1766 int btOpenCLSoftBodySolver::findSoftBodyIndex( const btSoftBody* const softBody )
1767 {
1768  for( int softBodyIndex = 0; softBodyIndex < m_softBodySet.size(); ++softBodyIndex )
1769  {
1770  btOpenCLAcceleratedSoftBodyInterface* softBodyInterface = m_softBodySet[softBodyIndex];
1771  if( softBodyInterface->getSoftBody() == softBody )
1772  return softBodyIndex;
1773  }
1774  return 1;
1775 }
1776 
1777 bool btOpenCLSoftBodySolver::checkInitialized()
1778 {
1779  if( !m_shadersInitialized )
1780  if( buildShaders() )
1781  m_shadersInitialized = true;
1782 
1783  return m_shadersInitialized;
1784 }
1785 
1786 bool btOpenCLSoftBodySolver::buildShaders()
1787 {
1788  if( m_shadersInitialized )
1789  return true;
1790 
1791  const char* additionalMacros="";
1792 
1793  // Ensure current kernels are released first
1794  releaseKernels();
1795 
1796  m_currentCLFunctions->clearKernelCompilationFailures();
1797 
1798  m_prepareLinksKernel = m_currentCLFunctions->compileCLKernelFromString( PrepareLinksCLString, "PrepareLinksKernel",additionalMacros,"OpenCLC10/PrepareLinks.cl" );
1799  m_updatePositionsFromVelocitiesKernel = m_currentCLFunctions->compileCLKernelFromString( UpdatePositionsFromVelocitiesCLString, "UpdatePositionsFromVelocitiesKernel" ,additionalMacros,"OpenCLC10/UpdatePositionsFromVelocities.cl");
1800  m_solvePositionsFromLinksKernel = m_currentCLFunctions->compileCLKernelFromString( SolvePositionsCLString, "SolvePositionsFromLinksKernel",additionalMacros,"OpenCLC10/SolvePositions.cl" );
1801  m_vSolveLinksKernel = m_currentCLFunctions->compileCLKernelFromString( VSolveLinksCLString, "VSolveLinksKernel" ,additionalMacros,"OpenCLC10/VSolveLinks.cl");
1802  m_updateVelocitiesFromPositionsWithVelocitiesKernel = m_currentCLFunctions->compileCLKernelFromString( UpdateNodesCLString, "updateVelocitiesFromPositionsWithVelocitiesKernel" ,additionalMacros,"OpenCLC10/UpdateNodes.cl");
1803  m_updateVelocitiesFromPositionsWithoutVelocitiesKernel = m_currentCLFunctions->compileCLKernelFromString( UpdatePositionsCLString, "updateVelocitiesFromPositionsWithoutVelocitiesKernel" ,additionalMacros,"OpenCLC10/UpdatePositions.cl");
1804  m_solveCollisionsAndUpdateVelocitiesKernel = m_currentCLFunctions->compileCLKernelFromString( SolveCollisionsAndUpdateVelocitiesCLString, "SolveCollisionsAndUpdateVelocitiesKernel" ,additionalMacros,"OpenCLC10/SolveCollisionsAndUpdateVelocities.cl");
1805  m_integrateKernel = m_currentCLFunctions->compileCLKernelFromString( IntegrateCLString, "IntegrateKernel" ,additionalMacros,"OpenCLC10/Integrate.cl");
1806  m_applyForcesKernel = m_currentCLFunctions->compileCLKernelFromString( ApplyForcesCLString, "ApplyForcesKernel" ,additionalMacros,"OpenCLC10/ApplyForces.cl");
1807  m_updateFixedVertexPositionsKernel = m_currentCLFunctions->compileCLKernelFromString( UpdateFixedVertexPositionsCLString, "UpdateFixedVertexPositions" , additionalMacros, "OpenCLC10/UpdateFixedVertexPositions.cl");
1808 
1809  // TODO: Rename to UpdateSoftBodies
1810  m_resetNormalsAndAreasKernel = m_currentCLFunctions->compileCLKernelFromString( UpdateNormalsCLString, "ResetNormalsAndAreasKernel" ,additionalMacros,"OpenCLC10/UpdateNormals.cl");
1811  m_normalizeNormalsAndAreasKernel = m_currentCLFunctions->compileCLKernelFromString( UpdateNormalsCLString, "NormalizeNormalsAndAreasKernel" ,additionalMacros,"OpenCLC10/UpdateNormals.cl");
1812  m_updateSoftBodiesKernel = m_currentCLFunctions->compileCLKernelFromString( UpdateNormalsCLString, "UpdateSoftBodiesKernel" ,additionalMacros,"OpenCLC10/UpdateNormals.cl");
1813 
1814 
1815  if( m_currentCLFunctions->getKernelCompilationFailures()==0 )
1816  m_shadersInitialized = true;
1817 
1818  return m_shadersInitialized;
1819 }
1820 
Generated on Thu Oct 3 2013 18:25:07 for Bullet Collision Detection & Physics Library by   doxygen 1.8.4