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 
18 #include "vectormath/vmInclude.h"
19 #include <stdio.h> //@todo: remove the debugging printf at some stage
26 #include "LinearMath/btQuickprof.h"
27 #include <limits.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 
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 
93 {
94 
95 }
96 
98 {
99  return m_onGPU;
100 }
101 
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 
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 
178 {
179 }
180 
182 {
183  Vectormath::Aos::Vector3 outVec( vec.getX(), vec.getY(), vec.getZ() );
184  return outVec;
185 }
186 
189 {
190  int previousSize = m_links.size();
191  int newSize = previousSize + numLinks;
192 
194 
195  // Resize the link addresses array as well
196  m_linkAddresses.resize( newSize );
197 }
198 
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 
211 {
212  return m_onGPU;
213 }
214 
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();
226 
227  if( success ) {
228  m_onGPU = true;
229  }
230 
231  return success;
232 }
233 
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();
245 
246  if( success ) {
247  m_onGPU = false;
248  }
249 
250  return success;
251 }
252 
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
344  btAlignedObjectArray<float> m_linkStrength_Backup(m_linkStrength);
345  btAlignedObjectArray<float> m_linksMassLSC_Backup(m_linksMassLSC);
346  btAlignedObjectArray<float> m_linksRestLengthSquared_Backup(m_linksRestLengthSquared);
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 
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 
401 {
402 }
403 
406 {
407  int previousSize = getNumTriangles();
408  int newSize = previousSize + numTriangles;
409 
411 
412  // Resize the link addresses array as well
413  m_triangleAddresses.resize( newSize );
414 }
415 
418 {
419  btSoftBodyTriangleData::setTriangleAt( triangle, triangleIndex );
420 
421  m_triangleAddresses[triangleIndex] = triangleIndex;
422 }
423 
425 {
426  return m_onGPU;
427 }
428 
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 
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 
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());
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
565  btAlignedObjectArray<float> m_area_Backup(m_area);
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 
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
631 
632  m_shadersInitialized = false;
633 
637  m_integrateKernel = 0;
650 }
651 
653 {
654  releaseKernels();
655 }
656 
658 {
674 
675  m_shadersInitialized = false;
676 }
677 
679 {
680 
681  // Move the vertex data back to the host first
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  {
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 
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();
718  getLinkData().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 ];
730 
731  // Create SoftBody that will store the information within the solver
733  m_softBodySet.push_back( newSoftBody );
734 
741  // Simple init values. Actually we'll put 0 and -1 into them at the appropriate time
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());
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 
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
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 
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
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 
888 
889  // generate batches
892 
893  // Build the shaders to match the batching parameters
894  buildShaders();
895  }
896 }
897 
898 
900 {
901  // TODO: Consider setting link data to "changed" here
902  return m_linkData;
903 }
904 
906 {
907  // TODO: Consider setting vertex data to "changed" here
908  return m_vertexData;
909 }
910 
912 {
913  // TODO: Consider setting triangle data to "changed" here
914  return m_triangleData;
915 }
916 
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 
938 {
939  cl_int ciErrNum;
940 
941  ciErrNum = clSetKernelArg(m_normalizeNormalsAndAreasKernel, 0, sizeof(int),(void*) &numVertices);
945  size_t numWorkItems = m_defaultWorkGroupSize*((numVertices + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize);
946  if (numWorkItems)
947  {
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);
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 
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
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 
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 
1025 {
1026  // Ensure data is on accelerator
1030 
1031  cl_int ciErrNum ;
1032  int numVerts = m_vertexData.getNumVertices();
1033  ciErrNum = clSetKernelArg(m_updateFixedVertexPositionsKernel, 0, sizeof(int), &numVerts);
1037 
1039  if (numWorkItems)
1040  {
1042  if( ciErrNum != CL_SUCCESS )
1043  {
1044  btAssert( 0 && "enqueueNDRangeKernel(m_updateFixedVertexPositionsKernel)");
1045  }
1046  }
1047 
1048 }
1049 
1051 {
1052  // Ensure data is on accelerator
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);
1077 
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 
1094 {
1095  // Ensure data is on accelerator
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);
1107 
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 
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 
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 
1145 {
1146 
1147  using namespace Vectormath::Aos;
1148 
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  {
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 
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 
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 
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 
1235 {
1236 
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 
1252 
1253 
1254  // Ensure data is on accelerator
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 
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  {
1282  } else {
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
1309 {
1310  cl_int ciErrNum;
1311  int numLinks = m_linkData.getNumLinks();
1312  ciErrNum = clSetKernelArg(m_prepareLinksKernel,0, sizeof(int), &numLinks);
1318 
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 
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);
1337 
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);
1359 
1360  size_t numWorkItems = m_defaultWorkGroupSize*((numLinks + (m_defaultWorkGroupSize-1)) / m_defaultWorkGroupSize);
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);
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 
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);
1404 
1407  if( ciErrNum != CL_SUCCESS )
1408  {
1409  btAssert( 0 && "enqueueNDRangeKernel(m_updateVelocitiesFromPositionsWithVelocitiesKernel)");
1410  }
1411 
1412 
1413 } // updateVelocitiesFromPositionsWithVelocities
1414 
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);
1427 
1430  if( ciErrNum != CL_SUCCESS )
1431  {
1432  btAssert( 0 && "enqueueNDRangeKernel(m_updateVelocitiesFromPositionsWithoutVelocitiesKernel)");
1433  }
1434 
1435 } // updateVelocitiesFromPositionsWithoutVelocities
1436 
1437 
1438 
1440 {
1441  // Copy kernel parameters to GPU
1445  m_clPerClothCollisionObjects.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);
1461 
1463  if (numWorkItems)
1464  {
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 
1481 {
1482 
1483  btSoftBodySolver *solver = softBody->getSoftBodySolver();
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;
1602  msg = "There is no successfully built executable for program.";
1603  break;
1605  msg = "kernel_name is not found in program.";
1606  break;
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 
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
1643  m_clPerClothCollisionObjects.changedOnCPU();
1645 
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 
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 )
1663 
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 
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 
1705 {
1706 
1707 }
1708 
1709 // Add the collision object to the set to deal with for a particular soft body
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 
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 
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 
1778 {
1779  if( !m_shadersInitialized )
1780  if( buildShaders() )
1781  m_shadersInitialized = true;
1782 
1783  return m_shadersInitialized;
1784 }
1785 
1787 {
1788  if( m_shadersInitialized )
1789  return true;
1790 
1791  const char* additionalMacros="";
1792 
1793  // Ensure current kernels are released first
1794  releaseKernels();
1795 
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 
1816  m_shadersInitialized = true;
1817 
1818  return m_shadersInitialized;
1819 }
1820 
static T sum(const btAlignedObjectArray< T > &items)
struct _cl_device_id * cl_device_id
Definition: cl.h:40
btOpenCLBuffer< Vectormath::Aos::Point3 > m_clAnchorPosition
virtual void setTriangleAt(const TriangleDescription &triangle, int triangleIndex)
btOpenCLBuffer< Vectormath::Aos::Point3 > m_clVertexPosition
btAlignedObjectArray< int > m_linkAddresses
Link addressing information for each cloth.
btScalar length(const btQuaternion &q)
Return the length of a quaternion.
Definition: btQuaternion.h:835
void createVertices(int numVertices, int clothIdentifier, int maxVertices=0)
Create numVertices new vertices for cloth clothIdentifier maxVertices allows a buffer zone of extra v...
btAlignedObjectArray< float > m_perClothDragFactor
Drag parameter for wind effect on cloth.
void push_back(const T &_Val)
virtual void solveCollisionsAndUpdateVelocities(float isolverdt)
struct _cl_context * cl_context
Definition: cl.h:41
int getShapeType() const
btOpenCLBuffer< float > m_clPerClothLiftFactor
virtual void setNumberOfPositionIterations(int iterations)
Set the number of velocity constraint solver iterations this solver uses.
#define CL_INVALID_KERNEL_NAME
Definition: cl.h:121
CL_API_ENTRY cl_kernel CL_API_CALL clCreateKernel(cl_program, const char *, cl_int *) CL_API_SUFFIX__VERSION_1_0
Definition: MiniCL.cpp:533
Config m_cfg
Definition: btSoftBody.h:648
static const char m_clVertexInverseMass(queue, ctx,&m_vertexInverseMass, false)
btOpenCLBuffer< Vectormath::Aos::Vector3 > m_clNormal
Vectormath::Aos::Vector3 linearVelocity
virtual void createTriangles(int numTriangles)
btOpenCLBuffer< Vectormath::Aos::Vector3 > m_clLinksCLength
void setPosition(const Vectormath::Aos::Point3 &position)
virtual void applyForces(float solverdt)
virtual bool hasVertexPositions() const
virtual void setLinkAt(const LinkDescription &link, int linkIndex)
Insert the link described into the correct data structures assuming space has already been allocated ...
btOpenCLBuffer< Vectormath::Aos::Point3 > m_clVertexPreviousPosition
virtual bool onAccelerator()
Return true if data is on the accelerator.
btAlignedObjectArray< Vectormath::Aos::Vector3 > m_normal
CL_API_ENTRY cl_program CL_API_CALL clCreateProgramWithSource(cl_context, cl_uint, const char **, const size_t *, cl_int *) CL_API_SUFFIX__VERSION_1_0
Definition: MiniCL.cpp:391
static Vectormath::Aos::Vector3 toVector3(const btVector3 &vec)
btAlignedObjectArray< TriangleNodeSet > m_vertexIndices
virtual void predictMotion(float solverdt)
Predict motion of soft bodies into next timestep.
virtual bool onAccelerator()
Return true if data is on the accelerator.
btOpenCLBuffer< float > m_clPerClothFriction
#define CL_PROGRAM_BUILD_LOG
Definition: cl.h:349
virtual BufferTypes getBufferType() const =0
Return the type of the vertex buffer descriptor.
float computeTriangleArea(const Vectormath::Aos::Point3 &vertex0, const Vectormath::Aos::Point3 &vertex1, const Vectormath::Aos::Point3 &vertex2)
btOpenCLBuffer< Vectormath::Aos::Vector3 > m_clVertexForceAccumulator
float & getRestLengthSquared(int linkIndex)
Return reference to rest length squared for link linkIndex as stored on the host. ...
#define RELEASE_CL_KERNEL(kernelName)
void solveLinksForVelocity(int startLink, int numLinks, float kst)
#define CL_PROGRAM_DEVICES
Definition: cl.h:341
btOpenCLBuffer< float > m_clLinkStrength
virtual bool moveToAccelerator()
Move data from host memory to the accelerator.
void generateBatches()
Generate (and later update) the batching for the entire triangle set.
The btCapsuleShape represents a capsule around the Y axis, there is also the btCapsuleShapeX aligned ...
btAlignedObjectArray< CollisionObjectIndices > m_perClothCollisionObjects
Collision shape details: pair of index of first collision shape for the cloth and number of collision...
int & getTriangleCount(int vertexIndex)
Get access to the array of how many triangles touch each vertex.
virtual void copySoftBodyToVertexBuffer(const btSoftBody *const softBody, btVertexBufferDescriptor *vertexBuffer)
Output current computed vertex data to the vertex buffers for all cloths in the solver.
void resetNormalsAndAreas(int numVertices)
#define btAssert(x)
Definition: btScalar.h:101
virtual bool moveFromAccelerator(bool bCopy=false, bool bCopyMinimum=true)
Move data to host memory from the accelerator if bCopy is false.
btSoftBodyLinkDataOpenCL m_linkData
int32_t cl_int
Definition: cl_platform.h:70
The btCollisionShape class provides an interface for collision shapes that can be shared among btColl...
virtual cl_kernel compileCLKernelFromString(const char *kernelSource, const char *kernelName, const char *additionalMacros, const char *srcFileNameForCaching)
Compile a compute shader kernel from a string and return the appropriate cl_kernel object...
btOpenCLBuffer< float > m_clLinksLengthRatio
static const char * PrepareLinksCLString
SoftBody class to maintain information about a soft body instance within a solver.
virtual bool moveFromAccelerator()
Move data from host memory from the accelerator.
btOpenCLBuffer< float > m_clLinksRestLength
float & getLinearStiffnessCoefficient(int linkIndex)
Return reference to linear stiffness coefficient for link linkIndex as stored on the host...
#define BT_DEFAULT_WORKGROUPSIZE
LinkNodePair & getVertexPair(int linkIndex)
Return reference to the vertex index pair for link linkIndex as stored on the host.
void ApplyClampedForce(float solverdt, const Vectormath::Aos::Vector3 &force, const Vectormath::Aos::Vector3 &vertexVelocity, float inverseMass, Vectormath::Aos::Vector3 &vertexForce)
static const char m_clVertexPosition(queue, ctx,&m_vertexPosition, false)
CL_API_ENTRY cl_int CL_API_CALL clEnqueueNDRangeKernel(cl_command_queue, cl_kernel, cl_uint, const size_t *, const size_t *, const size_t *, cl_uint, const cl_event *, cl_event *) CL_API_SUFFIX__VERSION_1_0
Definition: MiniCL.cpp:429
Start and length values for computation batches over link data.
virtual bool checkInitialized()
Ensure that this solver is initialized.
Class describing a vertex for input into the system.
btVector3 getColumn(int i) const
Get a column of the matrix as a vector.
Definition: btMatrix3x3.h:134
void normalizeNormalsAndAreas(int numVertices)
btAlignedObjectArray< Vectormath::Aos::Vector3 > m_linksCLength
int getKernelCompilationFailures() const
cl_kernel m_updateVelocitiesFromPositionsWithVelocitiesKernel
virtual void setTriangleAt(const btSoftBodyTriangleData::TriangleDescription &triangle, int triangleIndex)
Insert the link described into the correct data structures assuming space has already been allocated ...
btSoftBodyTriangleDataOpenCL(cl_command_queue queue, cl_context ctx)
virtual bool moveToAccelerator()
Move data from host memory to the accelerator.
btOpenCLBuffer< float > m_clLinksMassLSC
CL_API_ENTRY cl_int CL_API_CALL clGetProgramBuildInfo(cl_program, cl_device_id, cl_program_build_info, size_t, void *, size_t *) CL_API_SUFFIX__VERSION_1_0
Definition: MiniCL.cpp:377
tNodeArray m_nodes
Definition: btSoftBody.h:654
static const char * UpdatePositionsFromVelocitiesCLString
tLinkArray m_links
Definition: btSoftBody.h:655
btSoftBodyTriangleDataOpenCL m_triangleData
int findSoftBodyIndex(const btSoftBody *const softBody)
btAlignedObjectArray< btOpenCLAcceleratedSoftBodyInterface * > m_softBodySet
Cloths owned by this solver.
static float4 normalize(const float4 &a)
const btScalar & getZ() const
Return the z value.
Definition: btVector3.h:565
btAlignedObjectArray< Vectormath::Aos::Vector3 > m_perClothWindVelocity
Wind velocity to be applied normal to all non-static vertices in the solver.
btOpenCLBuffer< CollisionShapeDescription > m_clCollisionObjectDetails
void clear()
clear the array, deallocated memory. Generally it is better to use array.resize(0), to reduce performance overhead of run-time memory (de)allocations.
btOpenCLBuffer< Vectormath::Aos::Vector3 > m_clPerClothAcceleration
btAlignedObjectArray< float > m_linkStrength
Transform3 & setCol(int col, const Vector3 &vec)
float & getMassLSC(int linkIndex)
Return reference to the MassLSC value for link linkIndex as stored on the host.
static const char m_clVertexNormal(queue, ctx,&m_vertexNormal, false)
virtual void copyBackToSoftBodies(bool bMove=true)
Copy necessary data back to the original soft body source objects.
#define CL_SUCCESS
Definition: cl.h:91
float lengthSqr(const Vector3 &vec)
Definition: neon/vec_aos.h:447
btSoftBodyWorldInfo * getWorldInfo()
Definition: btSoftBody.h:697
btAlignedObjectArray< Vectormath::Aos::Vector3 > m_perClothAcceleration
Acceleration value to be applied to all non-static vertices in the solver.
static const char m_clVertexTriangleCount(queue, ctx,&m_vertexTriangleCount, false)
btAlignedObjectArray< AnchorNodeInfoCL > m_anchorNodeInfoArray
Vectormath::Aos::Transform3 shapeTransform
btOpenCLBuffer< btSoftBodyTriangleData::TriangleNodeSet > m_clVertexIndices
btOpenCLSoftBodySolver(cl_command_queue queue, cl_context ctx, bool bUpdateAchchoredNodePos=false)
int size() const
return the number of elements in the array
btSoftBodyLinkDataOpenCL(cl_command_queue queue, cl_context ctx)
btVector3 & getOrigin()
Return the origin vector translation.
Definition: btTransform.h:117
virtual btSoftBodyLinkData & getLinkData()
const btVector3 & getWindVelocity()
Return the wind velocity for interaction with the air.
Entry in the collision shape array.
virtual SolverTypes getSolverType() const =0
Return the type of the solver.
#define CL_INVALID_KERNEL_DEFINITION
Definition: cl.h:122
struct _cl_kernel * cl_kernel
Definition: cl.h:45
btAlignedObjectArray< Vectormath::Aos::Point3 > & getVertexPositions()
btAlignedObjectArray< LinkNodePair > m_links
Class representing a link as a set of three indices into the vertex array.
btSoftBodyVertexDataOpenCL(cl_command_queue queue, cl_context ctx)
btSoftBodySolver * getSoftBodySolver()
Definition: btSoftBody.h:913
const btVector3 & getAngularVelocity() const
Definition: btRigidBody.h:359
btAlignedObjectArray< int > m_anchorIndex
const btScalar & getY() const
Return the y value.
Definition: btVector3.h:563
const btTransform & getWorldTransform() const
static const char m_clLinkStrength(queue, ctx,&m_linkStrength, false)
btMatrix3x3 & getBasis()
Return the basis matrix for the rotation.
Definition: btTransform.h:112
const btScalar & getX() const
Return the x value.
Definition: btVector3.h:561
static const char m_clLinksRestLengthSquared(queue, ctx,&m_linksRestLengthSquared, false)
tAnchorArray m_anchors
Definition: btSoftBody.h:658
The btRigidBody is the main class for rigid body objects.
Definition: btRigidBody.h:59
virtual void integrate(float solverdt)
Integrate motion on the solver.
void updateBounds(const btVector3 &lowerBound, const btVector3 &upperBound)
Update the bounds in the btSoftBody object.
float & getRestLength(int linkIndex)
Return reference to the rest length of link linkIndex as stored on the host.
btAlignedObjectArray< float > m_linksRestLengthSquared
btAlignedObjectArray< float > m_linksLengthRatio
btAlignedObjectArray< BatchPair > m_batchStartLengths
Start and length values for computation batches over link data.
btAlignedObjectArray< float > m_linksMaterialLinearStiffnessCoefficient
virtual void processCollision(btSoftBody *, const btCollisionObjectWrapper *)
struct _cl_mem * cl_mem
Definition: cl.h:43
btAlignedObjectArray< float > m_perClothDampingFactor
Velocity damping factor.
CL_API_ENTRY cl_int CL_API_CALL clSetKernelArg(cl_kernel, cl_uint, size_t, const void *) CL_API_SUFFIX__VERSION_1_0
Definition: MiniCL.cpp:493
btAlignedObjectArray< float > m_perClothVelocityCorrectionCoefficient
Velocity correction coefficient.
virtual int getVertexStride() const
Return the vertex stride in number of floats between vertices.
btAlignedObjectArray< float > m_perClothFriction
Friction coefficient for each cloth.
static const char m_clVertexArea(queue, ctx,&m_vertexArea, false)
#define CL_INVALID_PROGRAM
Definition: cl.h:119
const btCollisionShape * getCollisionShape() const
void updateVelocitiesFromPositionsWithVelocities(float isolverdt)
btOpenCLBuffer< Vectormath::Aos::Vector3 > m_clVertexVelocity
void generateBatches()
Generate (and later update) the batching for the entire link set.
btVector3 can be used to represent 3D points and vectors.
Definition: btVector3.h:83
void prepareCollisionConstraints()
Sort the collision object details array and generate indexing into it for the per-cloth collision obj...
virtual int getVertexOffset() const
Return the vertex offset in floats from the base pointer.
Vectormath::Aos::Vector3 ProjectOnAxis(const Vectormath::Aos::Vector3 &v, const Vectormath::Aos::Vector3 &a)
btOpenCLBuffer< float > m_clLinksMaterialLinearStiffnessCoefficient
Class describing a link for input into the system.
#define BT_PROFILE(name)
Definition: btQuickprof.h:191
The btTransform class supports rigid transforms with only translation and rotation and no scaling/she...
Definition: btTransform.h:34
void clearKernelCompilationFailures()
static const char m_clLinksRestLength(queue, ctx,&m_linksRestLength, false)
btOpenCLBuffer< float > m_clPerClothMediumDensity
Wrapper for vertex data information.
btOpenCLBuffer< float > m_clLinksRestLengthSquared
CL_API_ENTRY cl_int CL_API_CALL clGetProgramInfo(cl_program, cl_program_info, size_t, void *, size_t *) CL_API_SUFFIX__VERSION_1_0
Definition: MiniCL.cpp:758
btOpenCLBuffer< float > m_clPerClothDampingFactor
virtual void solveLinksForPosition(int startLink, int numLinks, float kst, float ti)
void executeUpdateSoftBodies(int firstTriangle, int numTriangles)
Vectormath::Aos::Vector3 & getNormal(int vertexIndex)
Return a reference to the normal of vertex vertexIndex as stored on the host.
virtual void updateConstants(float timeStep)
btOpenCLAcceleratedSoftBodyInterface * findSoftBodyInterface(const btSoftBody *const softBody)
#define CL_INVALID_VALUE
Definition: cl.h:105
void resize(int newsize, const T &fillData=T())
btAlignedObjectArray< float > m_perClothLiftFactor
Lift parameter for wind effect on cloth.
void updateVelocitiesFromPositionsWithoutVelocities(float isolverdt)
Vectormath::Aos::Vector3 angularVelocity
btAlignedObjectArray< float > m_linksMassLSC
float & getInverseMass(int vertexIndex)
Return a reference to the inverse mass of vertex vertexIndex as stored on the host.
int findLinearSearch(const T &key) const
btAlignedObjectArray< btSomePair > m_batchStartLengths
virtual int getNormalOffset() const
Return the vertex offset in floats from the base pointer.
Vectormath::Aos::Point3 & getPosition(int vertexIndex)
Return a reference to the position of vertex vertexIndex as stored on the host.
btScalar getFriction() const
virtual float * getBasePointer() const
Return the base pointer in memory to the first vertex.
virtual btSoftBodyVertexData & getVertexData()
static float4 cross(const float4 &a, const float4 &b)
btOpenCLBuffer< Vectormath::Aos::Vector3 > m_clVertexNormal
struct _cl_command_queue * cl_command_queue
Definition: cl.h:42
static const char m_clLinksMaterialLinearStiffnessCoefficient(queue, ctx,&m_linksMaterialLinearStiffnessCoefficient, false)
btOpenCLBuffer< Vectormath::Aos::Vector3 > m_clPerClothWindVelocity
btAlignedObjectArray< float > m_area
static Vectormath::Aos::Transform3 toTransform3(const btTransform &transform)
static const char m_clLinksMassLSC(queue, ctx,&m_linksMassLSC, false)
float getTimeScale()
Return the timescale that the simulation is using.
btAlignedObjectArray< float > m_linksRestLength
#define CL_OUT_OF_HOST_MEMORY
Definition: cl.h:97
btAlignedObjectArray< CollisionShapeDescription > m_collisionObjectDetails
Collision shapes being passed across to the cloths in this solver.
void updatePositionsFromVelocities(float solverdt)
const btVector3 & getLinearVelocity() const
Definition: btRigidBody.h:356
virtual void optimize(btAlignedObjectArray< btSoftBody * > &softBodies, bool forceUpdate=false)
Optimize soft bodies in this solver.
btScalar air_density
Definition: btSoftBody.h:45
cl_kernel m_solveCollisionsAndUpdateVelocitiesKernel
virtual void updateSoftBodies()
Perform necessary per-step updates of soft bodies such as recomputing normals and bounding boxes...
const TriangleNodeSet & getVertexSet(int triangleIndex)
Return the vertex index set for triangle triangleIndex as stored on the host.
static const char m_clVertexForceAccumulator(queue, ctx,&m_vertexForceAccumulator, false)
struct _cl_program * cl_program
Definition: cl.h:44
btAlignedObjectArray< int > m_triangleAddresses
Link addressing information for each cloth.
btAlignedObjectArray< float > m_perClothMediumDensity
Density of the medium in which each cloth sits.
float dot(const Quat &quat0, const Quat &quat1)
virtual bool onAccelerator()
Return true if data is on the accelerator.
btVector3 m_gravity
Definition: btSoftBody.h:52
bool m_updateSolverConstants
Variable to define whether we need to update solver constants on the next iteration.
btOpenCLBuffer< float > m_clPerClothVelocityCorrectionCoefficient
CL_API_ENTRY cl_int CL_API_CALL clBuildProgram(cl_program, cl_uint, const cl_device_id *, const char *, void(*pfn_notify)(cl_program, void *), void *) CL_API_SUFFIX__VERSION_1_0
Definition: MiniCL.cpp:575
btOpenCLBuffer< float > m_clPerClothDragFactor
The btSoftBody is an class to simulate cloth and volumetric soft bodies.
Definition: btSoftBody.h:71
virtual btSoftBodyTriangleData & getTriangleData()
static const char m_clVertexVelocity(queue, ctx,&m_vertexVelocity, false)
virtual void setLinkAt(const LinkDescription &link, int linkIndex)
Insert the link described into the correct data structures assuming space has already been allocated ...
btVector3 m_x
Definition: btSoftBody.h:223
btSoftBodyVertexDataOpenCL m_vertexData
virtual void createTriangles(int numTriangles)
Allocate enough space in all link-related arrays to fit numLinks links.
btAlignedObjectArray< Vectormath::Aos::Point3 > m_anchorPosition
void setVertexAt(const VertexDescription &vertex, int vertexIndex)
virtual bool moveFromAccelerator()
Move data from host memory from the accelerator.
virtual bool moveToAccelerator()
Move data from host memory to the accelerator.
#define CL_INVALID_PROGRAM_EXECUTABLE
Definition: cl.h:120
virtual void setNumberOfVelocityIterations(int iterations)
Set the number of velocity constraint solver iterations this solver uses.
btOpenCLBuffer< LinkNodePair > m_clLinks
cl_kernel m_updateVelocitiesFromPositionsWithoutVelocitiesKernel
static const char m_clClothIdentifier(queue, ctx,&m_clothIdentifier, false)
void quickSort(const L &CompareFunc)
virtual int getNormalStride() const
Return the vertex stride in number of floats between vertices.
virtual void solveConstraints(float solverdt)
Solve constraints for a set of soft bodies.
const btCollisionObject * getCollisionObject() const
virtual void createLinks(int numLinks)
Allocate enough space in all link-related arrays to fit numLinks links.
btOpenCLBuffer< CollisionObjectIndices > m_clPerClothCollisionObjects
virtual void createLinks(int numLinks)
Allocate enough space in all link-related arrays to fit numLinks links.
tFaceArray m_faces
Definition: btSoftBody.h:656
btOpenCLBuffer< int > m_clAnchorIndex
static const char m_clVertexPreviousPosition(queue, ctx,&m_vertexPreviousPosition, false)