## On this page
# api\_cuda\_sph
Last update: 17.07.2025
1\#pragma once
2
3// Includes ---------------------------------------------------------------------------------------
4\#include <rocky20/api/device/api\_cuda.hpp>
5\#include <rocky20/api/sph/api\_device\_sph.hpp>
6\#include <rocky20/sph/sph\_data.h>
7
8namespace rocky::cuda {
9
10template <typename Operator, typename Functor>
11__global__ void execute_sph_particles_function(
12 SDeviceModel* device_data,
13 SPHDeviceModel* sph_device,
14 double current_time,
15 Functor sph_particles_f,
16 void* udf_device_data)
17{
18 device::launch_sph_particles_function<CudaBackend, Operator>(
19 device_data, sph_device, current_time, sph_particles_f, udf_device_data);
20}
21
22template <typename Operator, typename Functor>
23__global__ void execute_sph_interactions_function(
24 SDeviceModel* device_data,
25 SPHDeviceModel* sph_device,
26 double current_time,
27 Functor sph_particle_interaction_f,
28 void* udf_device_data,
29 IterationType it_type)
30{
31 // SPH -> SPH
32 device::launch_sph_interactions_particles<CudaBackend, Operator>(
33 device_data,
34 sph_device,
35 current_time,
36 sph_particle_interaction_f,
37 udf_device_data,
38 it_type);
39}
40
41template <typename Operator, typename Functor>
42__global__ void execute_sph_interactions_triangle_function(
43 SDeviceModel* device_data,
44 SPHDeviceModel* sph_device,
45 double current_time,
46 Functor sph_triangle_interaction_f,
47 void* udf_device_data)
48{
49 // SPH -> triangles
50 device::launch_sph_interactions_triangles<CudaBackend, Operator>(
51 device_data, sph_device, current_time, sph_triangle_interaction_f, udf_device_data);
52}
53
54// Dispatchers ------------------------------------------------------------------------------------
55template <typename Operator = device::DefaultOperator<CudaBackend>, typename Functor>
56void dispatch_sph_elements_kernel(
57 [IRockyModel](structIRockyModel.xhtml)& model, Functor& sph_particles_f, void** udf_devices_data)
58{
59 SPHModel* sph_model = &(*model._impl->sph_model);
60 cuda_host& cu_hos = *model._impl;
61
62 if (sph_model->gen_sph.nprt_active <= 0)
63 return;
64
65 void* udf_device_data = nullptr;
66
67 // Launching user kernel forces
68 for (int i = 0; i < model.[get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)(); ++i)
69 {
70 CUDA_SET_DEVICE(model.[get\_device\_id](structIRockyModel.xhtml#ae8f8865c809a31988291f9180d7cda2b)(i));
71
72 SPHDeviceModel* sph_device = model._impl->sph_model->get_device_data(i);
73 SDeviceModel* device_data = model.get_device_data(i);
74 if (udf_devices_data != nullptr)
75 {
76 udf_device_data = udf_devices_data[i];
77 }
78 rocky::cuda::execute_sph_particles_function<Operator><<<CD_NBLOCKS, CD_NTHREADS>>>(
79 device_data,
80 sph_device,
81 cu_hos.iteration.current_time,
82 sph_particles_f,
83 udf_device_data);
84 }
85
86 // Syncing all GPUs
87 sync_devices(model);
88}
89
90template <typename Operator = device::DefaultElementsOperator<CudaBackend>, typename Functor>
91void dispatch_sph_interactions_elements_kernel(
92 [IRockyModel](structIRockyModel.xhtml)& model,
93 Functor& sph_particle_interaction_f,
94 void** udf_devices_data,
95 IterationType it_type = itAll)
96{
97 SPHModel* sph_model = &(*model._impl->sph_model);
98 cuda_host& cu_hos = *model._impl;
99
100 if (sph_model->gen_sph.nprt_active <= 0)
101 return;
102
103 void* udf_device_data = nullptr;
104
105 // Launching user kernel forces
106 for (int i = 0; i < model.[get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)(); ++i)
107 {
108 CUDA_SET_DEVICE(model.[get\_device\_id](structIRockyModel.xhtml#ae8f8865c809a31988291f9180d7cda2b)(i));
109
110 SPHDeviceModel* sph_device = model._impl->sph_model->get_device_data(i);
111 SDeviceModel* device_data = model.get_device_data(i);
112 if (udf_devices_data != nullptr)
113 {
114 udf_device_data = udf_devices_data[i];
115 }
116 rocky::cuda::execute_sph_interactions_function<Operator><<<CD_NBLOCKS, CD_NTHREADS>>>(
117 device_data,
118 sph_device,
119 cu_hos.iteration.current_time,
120 sph_particle_interaction_f,
121 udf_device_data,
122 it_type);
123 }
124
125 // Syncing all GPUs
126 sync_devices(model);
127}
128
129template <typename Operator = device::DefaultTrianglesOperator<CudaBackend>, typename Functor>
130void dispatch_sph_interactions_triangles_kernel(
131 [IRockyModel](structIRockyModel.xhtml)& model, Functor& sph_triangle_interaction_f, void** udf_devices_data)
132{
133 SPHModel* sph_model = &(*model._impl->sph_model);
134 cuda_host& cu_hos = *model._impl;
135
136 if (sph_model->gen_sph.nprt_active <= 0)
137 return;
138
139 void* udf_device_data = nullptr;
140
141 // Launching user kernel forces
142 for (int i = 0; i < model.[get\_number\_of\_devices](structIRockyModel.xhtml#a1fd1011aa809fe64647e9cacf1956a11)(); ++i)
143 {
144 CUDA_SET_DEVICE(model.[get\_device\_id](structIRockyModel.xhtml#ae8f8865c809a31988291f9180d7cda2b)(i));
145
146 SPHDeviceModel* sph_device = model._impl->sph_model->get_device_data(i);
147 SDeviceModel* device_data = model.get_device_data(i);
148 if (udf_devices_data != nullptr)
149 {
150 udf_device_data = udf_devices_data[i];
151 }
152 rocky::cuda::execute_sph_interactions_triangle_function<Operator>
153 <<<CD_NBLOCKS, CD_NTHREADS>>>(
154 device_data,
155 sph_device,
156 cu_hos.iteration.current_time,
157 sph_triangle_interaction_f,
158 udf_device_data);
159 }
160
161 // Syncing all GPUs
162 sync_devices(model);
163}
164
165} // 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\_device\_id](structIRockyModel.xhtml#ae8f8865c809a31988291f9180d7cda2b)
int get_device_id(int device_index) const
**Definition** rocky_api.cpp:164