Rocky Solver SDK 2024 R2

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& 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(); ++i)
69 {
70 CUDA_SET_DEVICE(model.get_device_id(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& 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(); ++i)
107 {
108 CUDA_SET_DEVICE(model.get_device_id(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& 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(); ++i)
143 {
144 CUDA_SET_DEVICE(model.get_device_id(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
Definition rocky_api.h:60
IRockyModel::get_number_of_devices
int get_number_of_devices() const
Definition rocky_api.cpp:159
IRockyModel::get_device_id
int get_device_id(int device_index) const
Definition rocky_api.cpp:164

Connect with Ansys