PeriDyno 1.0.0
Loading...
Searching...
No Matches
LinearElasticitySolver.cu
Go to the documentation of this file.
1#include "LinearElasticitySolver.h"
2#include "Node.h"
3#include "Matrix/MatrixFunc.h"
4
5#include "ParticleSystem/Module/Kernel.h"
6
7namespace dyno
8{
9 IMPLEMENT_TCLASS(LinearElasticitySolver, TDataType)
10
11 template<typename Real>
12 __device__ Real D_Weight(Real r, Real h)
13 {
14 CorrectedKernel<Real> kernSmooth;
15 return kernSmooth.WeightRR(r, 2*h);
16// h = h < EPSILON ? Real(1) : h;
17// return 1 / (h*h*h);
18 }
19
20
21 template <typename Real, typename Coord, typename Matrix, typename Bond>
22 __global__ void EM_PrecomputeShape(
23 DArray<Matrix> invK,
24 DArray<Coord> X,
25 DArrayList<Bond> bonds)
26 {
27 int pId = threadIdx.x + (blockIdx.x * blockDim.x);
28 if (pId >= invK.size()) return;
29
30 List<Bond>& bonds_i = bonds[pId];
31
32 Coord rest_i = X[pId];
33
34 int size_i = bonds_i.size();
35 Real maxDist = Real(0);
36 for (int ne = 0; ne < size_i; ne++)
37 {
38 maxDist = max(maxDist, bonds_i[ne].xi.norm());
39 }
40 maxDist = maxDist < EPSILON ? Real(1) : maxDist;
41 Real smoothingLength = maxDist;
42
43 Real total_weight = 0.0f;
44 Matrix mat_i = Matrix(0);
45 for (int ne = 0; ne < size_i; ne++)
46 {
47 Bond bond_ij = bonds_i[ne];
48
49 int j = bond_ij.idx;
50 Coord rest_j = X[j];
51 Real r = (rest_i - rest_j).norm();
52
53 if (r > EPSILON)
54 {
55 Real weight = D_Weight(r, smoothingLength);
56 Coord q = (rest_j - rest_i) / smoothingLength*sqrt(weight);
57
58 mat_i(0, 0) += q[0] * q[0]; mat_i(0, 1) += q[0] * q[1]; mat_i(0, 2) += q[0] * q[2];
59 mat_i(1, 0) += q[1] * q[0]; mat_i(1, 1) += q[1] * q[1]; mat_i(1, 2) += q[1] * q[2];
60 mat_i(2, 0) += q[2] * q[0]; mat_i(2, 1) += q[2] * q[1]; mat_i(2, 2) += q[2] * q[2];
61
62 total_weight += weight;
63 }
64 }
65
66 if (total_weight > EPSILON)
67 {
68 mat_i *= (1.0f / total_weight);
69 }
70
71 Matrix R(0), U(0), D(0), V(0);
72
73 polarDecomposition(mat_i, R, U, D, V);
74
75 if (mat_i.determinant() < EPSILON*smoothingLength)
76 {
77 Real threshold = 0.0001f*smoothingLength;
78 D(0, 0) = D(0, 0) > threshold ? 1.0 / D(0, 0) : 1.0;
79 D(1, 1) = D(1, 1) > threshold ? 1.0 / D(1, 1) : 1.0;
80 D(2, 2) = D(2, 2) > threshold ? 1.0 / D(2, 2) : 1.0;
81
82 mat_i = V * D*U.transpose();
83 }
84 else
85 mat_i = mat_i.inverse();
86
87 invK[pId] = mat_i;
88 }
89
90 template <typename Real, typename Coord, typename Matrix, typename Bond>
91 __global__ void EM_EnforceElasticity(
92 DArray<Coord> delta_position,
93 DArray<Real> weights,
94 DArray<Real> bulkCoefs,
95 DArray<Matrix> invK,
96 DArray<Coord> X,
97 DArray<Coord> Y,
98 DArrayList<Bond> bonds,
99 Real mu,
100 Real lambda)
101 {
102 int pId = threadIdx.x + (blockIdx.x * blockDim.x);
103 if (pId >= Y.size()) return;
104
105 List<Bond>& bonds_i = bonds[pId];
106 Coord rest_i = X[pId];
107
108
109 Coord cur_pos_i = Y[pId];
110 Coord accPos = Coord(0);
111 Real accA = Real(0);
112 Real bulk_i = bulkCoefs[pId];
113
114 Real maxDist = Real(0);
115 int size_i = bonds_i.size();
116 for (int ne = 0; ne < size_i; ne++)
117 {
118 maxDist = max(maxDist, bonds_i[ne].xi.norm());
119 }
120 maxDist = maxDist < EPSILON ? Real(1) : maxDist;
121 Real horizon = maxDist;
122
123
124 Real total_weight = 0.0f;
125 Matrix deform_i = Matrix(0.0f);
126 for (int ne = 0; ne < size_i; ne++)
127 {
128 Bond bond_ij = bonds_i[ne];
129 int j = bond_ij.idx;
130
131 Coord rest_j = X[j];
132 Real r = (rest_j - rest_i).norm();
133
134 if (r > EPSILON)
135 {
136 Real weight = D_Weight(r, horizon);
137
138 Coord p = (Y[j] - Y[pId]) / horizon;
139 Coord q = (rest_j - rest_i) / horizon*weight;
140
141 deform_i(0, 0) += p[0] * q[0]; deform_i(0, 1) += p[0] * q[1]; deform_i(0, 2) += p[0] * q[2];
142 deform_i(1, 0) += p[1] * q[0]; deform_i(1, 1) += p[1] * q[1]; deform_i(1, 2) += p[1] * q[2];
143 deform_i(2, 0) += p[2] * q[0]; deform_i(2, 1) += p[2] * q[1]; deform_i(2, 2) += p[2] * q[2];
144 total_weight += weight;
145 }
146 }
147
148
149 if (total_weight > EPSILON)
150 {
151 deform_i *= (1.0f / total_weight);
152 deform_i = deform_i * invK[pId];
153 }
154 else
155 {
156 total_weight = 1.0f;
157 }
158
159 //Check whether the reference shape is inverted, if yes, simply set K^{-1} to be an identity matrix
160 //Note other solutions are possible.
161 if ((deform_i.determinant()) < -0.001f)
162 {
163 deform_i = Matrix::identityMatrix();
164 }
165
166
167// //if (pId == 0)
168// {
169// Matrix mat_i = invK[pId];
170// printf("Mat %d: \n %f %f %f \n %f %f %f \n %f %f %f \n",
171// pId,
172// mat_i(0, 0), mat_i(0, 1), mat_i(0, 2),
173// mat_i(1, 0), mat_i(1, 1), mat_i(1, 2),
174// mat_i(2, 0), mat_i(2, 1), mat_i(2, 2));
175// }
176
177 for (int ne = 0; ne < size_i; ne++)
178 {
179 Bond bond_ij = bonds_i[ne];
180
181 int j = bond_ij.idx;
182 Coord rest_j = X[j];
183
184 Coord cur_pos_j = Y[j];
185 Real r = (rest_j - rest_i).norm();
186
187 if (r > 0.01f*horizon)
188 {
189 Real weight = D_Weight(r, horizon);
190
191 Coord rest_dir_ij = deform_i*(rest_i - rest_j);
192 Coord cur_dir_ij = cur_pos_i - cur_pos_j;
193
194 cur_dir_ij = cur_dir_ij.norm() > EPSILON ? cur_dir_ij.normalize() : Coord(0);
195 rest_dir_ij = rest_dir_ij.norm() > EPSILON ? rest_dir_ij.normalize() : Coord(0, 0, 0);
196
197 Real mu_ij = mu*bulk_i* D_Weight(r, horizon);
198 Coord mu_pos_ij = Y[j] + r*rest_dir_ij;
199 Coord mu_pos_ji = Y[pId] - r*rest_dir_ij;
200
201 Real lambda_ij = lambda*bulk_i*D_Weight(r, horizon);
202 Coord lambda_pos_ij = Y[j] + r*cur_dir_ij;
203 Coord lambda_pos_ji = Y[pId] - r*cur_dir_ij;
204
205 Coord delta_pos_ij = mu_ij*mu_pos_ij + lambda_ij*lambda_pos_ij;
206 Real delta_weight_ij = mu_ij + lambda_ij;
207
208 Coord delta_pos_ji = mu_ij*mu_pos_ji + lambda_ij*lambda_pos_ji;
209
210 accA += delta_weight_ij;
211 accPos += delta_pos_ij;
212
213
214 atomicAdd(&weights[j], delta_weight_ij);
215 atomicAdd(&delta_position[j][0], delta_pos_ji[0]);
216 atomicAdd(&delta_position[j][1], delta_pos_ji[1]);
217 atomicAdd(&delta_position[j][2], delta_pos_ji[2]);
218 }
219 }
220
221 atomicAdd(&weights[pId], accA);
222 atomicAdd(&delta_position[pId][0], accPos[0]);
223 atomicAdd(&delta_position[pId][1], accPos[1]);
224 atomicAdd(&delta_position[pId][2], accPos[2]);
225 }
226
227 template <typename Real, typename Coord>
228 __global__ void K_UpdatePosition(
229 DArray<Coord> position,
230 DArray<Coord> old_position,
231 DArray<Coord> delta_position,
232 DArray<Real> delta_weights)
233 {
234 int pId = threadIdx.x + (blockIdx.x * blockDim.x);
235 if (pId >= position.size()) return;
236
237 position[pId] = (old_position[pId] + delta_position[pId]) / (1.0+delta_weights[pId]);
238 }
239
240 template <typename Real, typename Coord>
241 __global__ void K_UpdateVelocity(
242 DArray<Coord> velArr,
243 DArray<Coord> prePos,
244 DArray<Coord> curPos,
245 Real dt)
246 {
247 int pId = threadIdx.x + (blockIdx.x * blockDim.x);
248 if (pId >= velArr.size()) return;
249
250 velArr[pId] += (curPos[pId] - prePos[pId]) / dt;
251 }
252
253 template<typename TDataType>
254 LinearElasticitySolver<TDataType>::LinearElasticitySolver()
255 : ConstraintModule()
256 {
257 this->inHorizon()->setValue(0.0125);
258 }
259
260
261 template<typename TDataType>
262 LinearElasticitySolver<TDataType>::~LinearElasticitySolver()
263 {
264 mWeights.clear();
265 mDisplacement.clear();
266 mInvK.clear();
267 mF.clear();
268 mPosBuf.clear();
269 }
270
271 template<typename TDataType>
272 void LinearElasticitySolver<TDataType>::enforceElasticity()
273 {
274 int num = this->inY()->size();
275 uint pDims = cudaGridSize(num, BLOCK_SIZE);
276
277 mDisplacement.reset();
278 mWeights.reset();
279
280 EM_EnforceElasticity << <pDims, BLOCK_SIZE >> > (
281 mDisplacement,
282 mWeights,
283 mBulkStiffness,
284 mInvK,
285 this->inX()->getData(),
286 this->inY()->getData(),
287 this->inBonds()->getData(),
288 this->varMu()->getData(),
289 this->varLambda()->getData());
290 cuSynchronize();
291
292 K_UpdatePosition << <pDims, BLOCK_SIZE >> > (
293 this->inY()->getData(),
294 mPosBuf,
295 mDisplacement,
296 mWeights);
297 cuSynchronize();
298 }
299
300 template<typename Real>
301 __global__ void EM_InitBulkStiffness(DArray<Real> stiffness)
302 {
303 int pId = threadIdx.x + (blockIdx.x * blockDim.x);
304 if (pId >= stiffness.size()) return;
305
306 stiffness[pId] = Real(1);
307 }
308
309 template<typename TDataType>
310 void LinearElasticitySolver<TDataType>::computeMaterialStiffness()
311 {
312 int num = this->inY()->size();
313
314 uint pDims = cudaGridSize(num, BLOCK_SIZE);
315 EM_InitBulkStiffness << <pDims, BLOCK_SIZE >> > (mBulkStiffness);
316 }
317
318
319 template<typename TDataType>
320 void LinearElasticitySolver<TDataType>::computeInverseK()
321 {
322 auto& restShapes = this->inBonds()->getData();
323 uint pDims = cudaGridSize(restShapes.size(), BLOCK_SIZE);
324
325 EM_PrecomputeShape <Real, Coord, Matrix, Bond> << <pDims, BLOCK_SIZE >> > (
326 mInvK,
327 this->inX()->getData(),
328 restShapes);
329 cuSynchronize();
330 }
331
332 template<typename TDataType>
333 void LinearElasticitySolver<TDataType>::solveElasticity()
334 {
335 //Save new positions
336 mPosBuf.assign(this->inY()->getData());
337
338 this->computeInverseK();
339
340 int itor = 0;
341 uint maxIterNum = this->varIterationNumber()->getData();
342 while (itor < maxIterNum) {
343 this->enforceElasticity();
344 itor++;
345 }
346
347 this->updateVelocity();
348 }
349
350 template<typename TDataType>
351 void LinearElasticitySolver<TDataType>::updateVelocity()
352 {
353 int num = this->inY()->size();
354 uint pDims = cudaGridSize(num, BLOCK_SIZE);
355
356 Real dt = this->inTimeStep()->getData();
357
358 K_UpdateVelocity << <pDims, BLOCK_SIZE >> > (
359 this->inVelocity()->getData(),
360 mPosBuf,
361 this->inY()->getData(),
362 dt);
363 cuSynchronize();
364 }
365
366
367 template<typename TDataType>
368 void LinearElasticitySolver<TDataType>::constrain()
369 {
370 this->solveElasticity();
371 }
372
373
374 template <typename Coord, typename Bond>
375 __global__ void K_UpdateRestShape(
376 DArrayList<Bond> shape,
377 DArrayList<int> nbr,
378 DArray<Coord> pos)
379 {
380 int pId = threadIdx.x + (blockIdx.x * blockDim.x);
381 if (pId >= pos.size()) return;
382
383 Bond np;
384
385 List<Bond>& rest_shape_i = shape[pId];
386 List<int>& list_id_i = nbr[pId];
387 int nbSize = list_id_i.size();
388 for (int ne = 0; ne < nbSize; ne++)
389 {
390 int j = list_id_i[ne];
391 np.index = j;
392 np.pos = pos[j];
393 np.weight = 1;
394
395 rest_shape_i.insert(np);
396 if (pId == j)
397 {
398 Bond np_0 = rest_shape_i[0];
399 rest_shape_i[0] = np;
400 rest_shape_i[ne] = np_0;
401 }
402 }
403 }
404
405 template<typename TDataType>
406 void LinearElasticitySolver<TDataType>::preprocess()
407 {
408 int num = this->inY()->size();
409
410 if (num == mInvK.size())
411 return;
412
413 mInvK.resize(num);
414 mWeights.resize(num);
415 mDisplacement.resize(num);
416
417 mF.resize(num);
418
419 mPosBuf.resize(num);
420 mBulkStiffness.resize(num);
421
422 this->computeMaterialStiffness();
423 }
424
425 DEFINE_CLASS(LinearElasticitySolver);
426}