## On this page
# api\_cuda\_contacts
Last update: 17.07.2025
1\#pragma once
2
3\#include <rocky20/api/device/api\_device\_contacts.hpp>
4\#include <rocky20/api/device/api\_cuda.hpp>
5\#include <rocky20/api/rocky\_api\_functions.h>
6\#include <rocky20/api/rocky\_api.h>
7
8namespace rocky::cuda
9{
10
11// =================================================================================================
12// Contacts Loop
13// =================================================================================================
14
15template <typename Callable>
16__global__ void execute_frictional_contacts_function(
17 SDeviceModel *device_data,
18 cu_real time_curr,
19 Callable particle_contact_f_ptr,
20 Callable triangle_contact_f_ptr,
21 void *udf_device_data
22)
23{
24 // Particles
25 device::execute_contacts_function<CudaBackend>(
26 device_data,
27 time_curr,
28 device_data->ind_cnt_p_fr,
29 device_data->n_contacts->particle_fr,
30 true,
31 particle_contact_f_ptr,
32 udf_device_data);
33
34 // Triangles
35 device::execute_contacts_function<CudaBackend>(
36 device_data,
37 time_curr,
38 device_data->ind_cnt_t_fr,
39 device_data->n_contacts->triangle_fr,
40 false,
41 triangle_contact_f_ptr,
42 udf_device_data);
43}
44
45template <typename Callable>
46__global__ void execute_finished_contacts_function(
47 SDeviceModel* device_data,
48 cu_real time_curr,
49 Callable particle_contact_f_ptr,
50 Callable triangle_contact_f_ptr,
51 void* udf_device_data
52)
53{
54 // particles -> particles
55 device::execute_finished_contacts_function<CudaBackend>(
56 device_data,
57 time_curr,
58 device_data->gen->ncont_p,
59 true,
60 particle_contact_f_ptr, udf_device_data
61 );
62
63 // particles -> triangles
64 device::execute_finished_contacts_function<CudaBackend>(
65 device_data,
66 time_curr,
67 device_data->gen->ncont_t,
68 false,
69 triangle_contact_f_ptr, udf_device_data
70 );
71}
72
73template <typename Callable>
74__global__ void execute_active_contacts_function(
75 SDeviceModel* device_data,
76 cu_real time_curr,
77 Callable particle_contact_f_ptr,
78 Callable triangle_contact_f_ptr,
79 void* udf_device_data
80)
81{
82 // particles -> particles
83 device::execute_active_contacts_function<CudaBackend>(
84 device_data,
85 time_curr,
86 device_data->gen->ncont_p,
87 true,
88 particle_contact_f_ptr,
89 udf_device_data);
90
91 // particles -> triangles
92 device::execute_active_contacts_function<CudaBackend>(
93 device_data,
94 time_curr,
95 device_data->gen->ncont_t,
96 false,
97 triangle_contact_f_ptr,
98 udf_device_data);
99}
100
101template <typename Callable>
102void dispatch_frictional_contacts_kernel(
103 [IRockyModel](structIRockyModel.xhtml) &model,
104 Callable h_particle_contact_f_ptr,
105 Callable h_triangle_contact_f_ptr,
106 void **udf_devices_data
107)
108{
109 if (model.[get\_number\_of\_active\_particles](structIRockyModel.xhtml#a38a15145f4f2bb47b79ab785f50d8277)() <= 0)
110 return;
111
112 for (int i = 0; i < model.[get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)(); i++)
113 {
114 void *udf_device_data = udf_devices_data[i];
115 SDeviceModel *device_data = model.get_device_data(i);
116
117 CUDA_SET_DEVICE(model.[get\_device\_id](structIRockyModel.xhtml#ae8f8865c809a31988291f9180d7cda2b)(i));
118 rocky::cuda::execute_frictional_contacts_function<<<CD_NBLOCKS, CD_NTHREADS>>>(
119 device_data,
120 model.[get\_current\_time](structIRockyModel.xhtml#a93ef35ae1eb9cf50ca77ab2070a769ab)(),
121 h_particle_contact_f_ptr,
122 h_triangle_contact_f_ptr,
123 udf_device_data);
124 }
125
126 //Sync all GPUs
127 sync_devices(model);
128}
129
130// frictional and adhesive contacts just finished
131template <typename Callable>
132void dispatch_finished_contacts_kernel(
133 [IRockyModel](structIRockyModel.xhtml)& model,
134 Callable h_particle_contact_f_ptr,
135 Callable h_triangle_contact_f_ptr,
136 void** udf_devices_data
137)
138{
139 if (model.[get\_number\_of\_active\_particles](structIRockyModel.xhtml#a38a15145f4f2bb47b79ab785f50d8277)() <= 0)
140 return;
141
142 for (int i = 0; i < model.[get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)(); i++)
143 {
144 void* udf_device_data = udf_devices_data[i];
145 SDeviceModel* device_data = model.get_device_data(i);
146
147 CUDA_SET_DEVICE(model.[get\_device\_id](structIRockyModel.xhtml#ae8f8865c809a31988291f9180d7cda2b)(i));
148 rocky::cuda::execute_finished_contacts_function <<<CD_NBLOCKS, CD_NTHREADS>>> (
149 device_data, model.[get\_current\_time](structIRockyModel.xhtml#a93ef35ae1eb9cf50ca77ab2070a769ab)(),
150 h_particle_contact_f_ptr, h_triangle_contact_f_ptr,
151 udf_device_data
152 );
153 }
154
155 sync_devices(model);
156}
157
158// frictional and adhesive contacts still active
159template <typename Callable>
160void dispatch_active_contacts_kernel(
161 [IRockyModel](structIRockyModel.xhtml)& model,
162 Callable h_particle_contact_f_ptr,
163 Callable h_triangle_contact_f_ptr,
164 void** udf_devices_data)
165{
166 if (model.[get\_number\_of\_active\_particles](structIRockyModel.xhtml#a38a15145f4f2bb47b79ab785f50d8277)() <= 0)
167 return;
168
169 for (int i = 0; i < model.[get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)(); i++)
170 {
171 void* udf_device_data = udf_devices_data[i];
172 SDeviceModel* device_data = model.get_device_data(i);
173
174 CUDA_SET_DEVICE(model.[get\_device\_id](structIRockyModel.xhtml#ae8f8865c809a31988291f9180d7cda2b)(i));
175 rocky::cuda::execute_active_contacts_function<<<CD_NBLOCKS, CD_NTHREADS>>>(
176 device_data,
177 model.[get\_current\_time](structIRockyModel.xhtml#a93ef35ae1eb9cf50ca77ab2070a769ab)(),
178 h_particle_contact_f_ptr,
179 h_triangle_contact_f_ptr,
180 udf_device_data);
181 }
182
183 sync_devices(model);
184}
185
186// Adhesive ----------------------------------------------------------------------------------------
187
188template <typename Callable>
189__global__ void execute_adhesive_contacts_function(
190 SDeviceModel *device_data,
191 cu_real time_curr,
192 Callable particle_contact_f_ptr,
193 Callable triangle_contact_f_ptr,
194 void *udf_device_data
195)
196{
197 // Particles
198 if (device_data->n_contacts->particle_ad > 0)
199 {
200 device::execute_contacts_function<CudaBackend>(
201 device_data,
202 time_curr,
203 device_data->ind_cnt_p_ad,
204 device_data->n_contacts->particle_ad,
205 true,
206 particle_contact_f_ptr,
207 udf_device_data);
208 }
209
210 // Triangles
211 if (device_data->n_contacts->triangle_ad > 0)
212 {
213 device::execute_contacts_function<CudaBackend>(
214 device_data,
215 time_curr,
216 device_data->ind_cnt_t_ad,
217 device_data->n_contacts->triangle_ad,
218 false,
219 triangle_contact_f_ptr,
220 udf_device_data);
221 }
222}
223
224template <typename Callable>
225void dispatch_adhesive_contacts_kernel(
226 [IRockyModel](structIRockyModel.xhtml) &model,
227 Callable h_particle_contact_f_ptr,
228 Callable h_triangle_contact_f_ptr,
229 void **udf_devices_data
230)
231{
232 if (model.[get\_number\_of\_active\_particles](structIRockyModel.xhtml#a38a15145f4f2bb47b79ab785f50d8277)() <= 0)
233 return;
234
235 for (int i = 0; i<model.[get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)(); i++)
236 {
237 void *udf_device_data = udf_devices_data[i];
238 SDeviceModel *device_data = model.get_device_data(i);
239
240 CUDA_SET_DEVICE(model.[get\_device\_id](structIRockyModel.xhtml#ae8f8865c809a31988291f9180d7cda2b)(i));
241 rocky::cuda::execute_adhesive_contacts_function<<<CD_NBLOCKS, CD_NTHREADS>>>(
242 device_data,
243 model.[get\_current\_time](structIRockyModel.xhtml#a93ef35ae1eb9cf50ca77ab2070a769ab)(),
244 h_particle_contact_f_ptr,
245 h_triangle_contact_f_ptr,
246 udf_device_data);
247 }
248
249 //Sync all GPUs
250 sync_devices(model);
251
252}
253
254// =================================================================================================
255// All Contacts Loop
256// =================================================================================================
257template <typename Callable>
258__global__ void execute_all_contacts_function(
259 SDeviceModel *device_data,
260 cu_real time_curr,
261 Callable particle_contact_f_ptr,
262 Callable triangle_contact_f_ptr,
263 void *udf_device_data
264)
265{
266 // Particles
267 device::execute_all_contacts_function<CudaBackend>(
268 device_data,
269 time_curr,
270 device_data->gen->ncont_p,
271 true,
272 particle_contact_f_ptr,
273 udf_device_data);
274
275 // Triangles
276 device::execute_all_contacts_function<CudaBackend>(
277 device_data,
278 time_curr,
279 device_data->gen->ncont_t,
280 false,
281 triangle_contact_f_ptr,
282 udf_device_data);
283}
284
285template <typename Callable>
286void dispatch_all_contacts_kernel(
287 [IRockyModel](structIRockyModel.xhtml) &model,
288 Callable h_particle_contact_f_ptr,
289 Callable h_triangle_contact_f_ptr,
290 void **udf_device_data
291)
292{
293 if (model.[get\_number\_of\_active\_particles](structIRockyModel.xhtml#a38a15145f4f2bb47b79ab785f50d8277)() <= 0)
294 return;
295
296 for (int i = 0; i<model.[get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)(); i++)
297 {
298 CUDA_SET_DEVICE(model.[get\_device\_id](structIRockyModel.xhtml#ae8f8865c809a31988291f9180d7cda2b)(i));
299 rocky::cuda::execute_all_contacts_function<<<CD_NBLOCKS, CD_NTHREADS>>>(
300 model.get_device_data(i),
301 model.[get\_current\_time](structIRockyModel.xhtml#a93ef35ae1eb9cf50ca77ab2070a769ab)(),
302 h_particle_contact_f_ptr,
303 h_triangle_contact_f_ptr,
304 udf_device_data[i]);
305 }
306
307 //Sync all GPUs
308 sync_devices(model);
309}
310
311// =================================================================================================
312// Contact forces
313// =================================================================================================
314template <typename CallableCollection>
315__global__ void execute_frictional_particle_contact_forces_function(
316 SDeviceModel* device_data,
317 cu_real time_curr,
318 const CallableCollection functions,
319 void* udf_device_data)
320{
321 device::api_compute_particle_contact_frictional_forces<CudaBackend>(
322 device_data, time_curr, functions, udf_device_data);
323}
324
325template <typename CallableCollection>
326__global__ void execute_frictional_triangle_contact_forces_function(
327 SDeviceModel* device_data,
328 cu_real time_curr,
329 const CallableCollection functions,
330 void* udf_device_data)
331{
332 device::api_compute_triangle_contact_frictional_forces<CudaBackend>(
333 device_data, time_curr, functions, udf_device_data);
334}
335
336
337template <typename CallableCollection>
338void dispatch_frictional_contact_forces_kernel(
339 [IRockyModel](structIRockyModel.xhtml)& model, const CallableCollection& function_ptrs, void** udf_devices_data)
340{
341 if (model.[get\_number\_of\_active\_particles](structIRockyModel.xhtml#a38a15145f4f2bb47b79ab785f50d8277)() <= 0)
342 return;
343
344 auto& solver_model = *model._impl;
345
346 solver_model.n_contacts.reset();
347 std::vector<SNumberContacts> n_contacts(model.[get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)());
348 for (int i = 0; i < model.[get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)(); i++)
349 {
350 CUDA_SET_DEVICE(model.[get\_device\_id](structIRockyModel.xhtml#ae8f8865c809a31988291f9180d7cda2b)(i));
351 CUDA_COPY_D2H(n_contacts.data() + i, solver_model.GetDeviceModel(i).n_contacts, 1);
352 solver_model.n_contacts.join(n_contacts[i]);
353 }
354
355 void* udf_device_data = nullptr;
356 for (int i = 0; i < model.[get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)(); i++)
357 {
358 if (udf_devices_data) udf_device_data = udf_devices_data[i];
359 SDeviceModel* device_data = model.get_device_data(i);
360
361 CUDA_SET_DEVICE(model.[get\_device\_id](structIRockyModel.xhtml#ae8f8865c809a31988291f9180d7cda2b)(i))
362 auto [blocks, threads] = CudaBackend::get_launch_params(n_contacts[i].particle_fr);
363 rocky::cuda::execute_frictional_particle_contact_forces_function<<<blocks, threads>>>(
364 device_data, model.[get\_current\_time](structIRockyModel.xhtml#a93ef35ae1eb9cf50ca77ab2070a769ab)(), function_ptrs, udf_device_data);
365 }
366
367 for (int i = 0; i < model.[get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)(); i++)
368 {
369 if (udf_devices_data) udf_device_data = udf_devices_data[i];
370 SDeviceModel* device_data = model.get_device_data(i);
371
372 CUDA_SET_DEVICE(model.[get\_device\_id](structIRockyModel.xhtml#ae8f8865c809a31988291f9180d7cda2b)(i))
373 auto [blocks, threads] = CudaBackend::get_launch_params(n_contacts[i].triangle_fr);
374 rocky::cuda::execute_frictional_triangle_contact_forces_function <<<blocks, threads>>> (
375 device_data, model.[get\_current\_time](structIRockyModel.xhtml#a93ef35ae1eb9cf50ca77ab2070a769ab)(), function_ptrs, udf_device_data);
376 }
377
378 sync_devices(model);
379}
380
381template <typename Callable>
382__global__ void execute_adhesive_contact_forces_function(
383 SDeviceModel *device_data,
384 cu_real time_curr,
385 Callable particle_contact_f_ptr,
386 Callable triangle_contact_f_ptr,
387 void *udf_device_data
388)
389{
390 // Particles
391 device::api_compute_particle_contact_adhesive_forces<CudaBackend>(
392 device_data, time_curr,
393 particle_contact_f_ptr, udf_device_data
394 );
395
396 // Triangles
397 device::api_compute_triangle_contact_adhesive_forces<CudaBackend>(
398 device_data, time_curr,
399 triangle_contact_f_ptr, udf_device_data
400 );
401}
402
403template <typename Callable>
404void dispatch_adhesive_contact_forces_kernel(
405 [IRockyModel](structIRockyModel.xhtml) &model,
406 Callable h_particle_contact_f_ptr,
407 Callable h_triangle_contact_f_ptr,
408 void **udf_devices_data
409)
410{
411 if (model.[get\_number\_of\_active\_particles](structIRockyModel.xhtml#a38a15145f4f2bb47b79ab785f50d8277)() <= 0)
412 return;
413
414 std::vector<SNumberContacts> n_contacts(model.[get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)());
415 for (int i = 0; i < model.[get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)(); i++)
416 {
417 CUDA_SET_DEVICE(model.[get\_device\_id](structIRockyModel.xhtml#ae8f8865c809a31988291f9180d7cda2b)(i))
418 CUDA_COPY_D2H(n_contacts.data() + i, model._impl->GetDeviceModel(i).n_contacts, 1)
419 }
420
421 for (int i = 0; i < model.[get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)(); i++)
422 {
423 void *udf_device_data = udf_devices_data[i];
424 SDeviceModel *device_data = model.get_device_data(i);
425
426 CUDA_SET_DEVICE(model.[get\_device\_id](structIRockyModel.xhtml#ae8f8865c809a31988291f9180d7cda2b)(i))
427 auto [blocks, threads] = CudaBackend::get_launch_params(
428 std::max(n_contacts[i].particle_ad, n_contacts[i].triangle_ad));
429 rocky::cuda::execute_adhesive_contact_forces_function<<<blocks, threads>>>(
430 device_data,
431 model.[get\_current\_time](structIRockyModel.xhtml#a93ef35ae1eb9cf50ca77ab2070a769ab)(),
432 h_particle_contact_f_ptr,
433 h_triangle_contact_f_ptr,
434 udf_device_data);
435 }
436
437 //Sync all GPUs
438 sync_devices(model);
439}
440
441// Boundary wear on particle-triangle contacts ----------------------------------------------------
442
443template <typename Callable>
444__global__ void execute_boundaries_wear_function_kernel(
445 SDeviceModel* device_data,
446 Callable triangle_wear_f_ptr,
447 void* udf_device_data,
448 cu_real time_curr)
449{
450 device::api_compute_boundaries_wear<CudaBackend>(
451 device_data, time_curr, triangle_wear_f_ptr, udf_device_data);
452}
453
454template <typename Callable>
455void dispatch_boundaries_wear_kernel(
456 [IRockyModel](structIRockyModel.xhtml)& model,
457 Callable triangle_wear_f_ptr,
458 void** udf_devices_data
459)
460{
461 for (int i = 0; i < model.[get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)(); i++)
462 {
463 void* udf_device_data = udf_devices_data[i];
464 SDeviceModel* device_data = model.get_device_data(i);
465
466 CUDA_SET_DEVICE(model.[get\_device\_id](structIRockyModel.xhtml#ae8f8865c809a31988291f9180d7cda2b)(i))
467 rocky::cuda::execute_boundaries_wear_function_kernel<<<CD_NBLOCKS, CD_NTHREADS>>>(
468 device_data,
469 triangle_wear_f_ptr,
470 udf_device_data,
471 model._impl->iteration.current_time);
472 }
473
474 sync_devices(model);
475}
476
477} // namespace rocky::cuda
[IRockyModel](structIRockyModel.xhtml)
**Definition** rocky_api.h:60
[IRockyModel::get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)
int get_number_of_devices() const
**Definition** rocky_api.cpp:159
[IRockyModel::get\_number\_of\_active\_particles](structIRockyModel.xhtml#a38a15145f4f2bb47b79ab785f50d8277)
int get_number_of_active_particles() const
**Definition** rocky_api.cpp:139
[IRockyModel::get\_current\_time](structIRockyModel.xhtml#a93ef35ae1eb9cf50ca77ab2070a769ab)
double get_current_time() const
**Definition** rocky_api.cpp:129
[IRockyModel::get\_device\_id](structIRockyModel.xhtml#ae8f8865c809a31988291f9180d7cda2b)
int get_device_id(int device_index) const
**Definition** rocky_api.cpp:164