api_omp_sph Last update: 16.07.2025 1// (C) 2024 ANSYS, Inc. Unauthorized use, distribution, or duplication is prohibited. 2#pragma once 3 4// Includes --------------------------------------------------------------------------------------- 5#include <rocky20/api/rocky_api.h> 6 7#include <rocky20/api/sph/api_device_sph.hpp> 8#include <rocky20/device/omp_backend.hpp> 9 10namespace rocky::omp { 11 12// Dispatchers ------------------------------------------------------------------------------------ 13template <typename Operator = device::DefaultOperator<OMPBackend>, typename Functor> 14inline void dispatch_sph_elements_kernel( 15 IRockyModel& model, Functor& omp_dispatch_particle_f, void* udf_data) 16{ 17 SPHModel* sph_model = &(*model._impl->sph_model); 18 SPHDeviceModel* sph_device = model._impl->sph_model->get_device_data(0); 19 SDeviceModel* device_data = model.get_device_data(0); 20 21 if (sph_model->gen_sph.nprt_active <= 0) 22 return; 23 24#ifndef GPU_KERNEL_COMPILING_STAGE 25 #pragma omp parallel num_threads(model.get_number_of_processes()) 26 { 27 device::launch_sph_particles_function<OMPBackend, Operator>( 28 device_data, sph_device, model.get_current_time(), omp_dispatch_particle_f, udf_data); 29 } 30#endif 31} 32 33template <typename Operator = device::DefaultElementsOperator<OMPBackend>, typename Functor> 34inline void dispatch_sph_interactions_elements_kernel( 35 IRockyModel& model, 36 Functor& sph_particle_interaction_f, 37 void* udf_data, 38 IterationType it_type = itAll) 39{ 40 cuda_host& solver_model = *model._impl; 41 SPHModel* sph_model = &(*model._impl->sph_model); 42 float current_time = model.get_current_time(); 43 if (sph_model->gen_sph.nprt_active <= 0) 44 return; 45 46#ifndef GPU_KERNEL_COMPILING_STAGE 47 #pragma omp parallel num_threads(model.get_number_of_processes()) 48 { 49 // SPH -> SPH 50 if (sph_model->gen_sph.use_particles_neighbors_list) 51 { 52 if (solver_model.gen.periodic_domain.enabled()) 53 { 54 SPHExplicitHandler<OMPBackend, true> handler { solver_model.GetDeviceData(0), 55 sph_model->get_device_data(0), 56 current_time, 57 it_type }; 58 device::launch_sph_interactions_particles< 59 OMPBackend, 60 Operator, 61 decltype(sph_particle_interaction_f)>( 62 solver_model.GetDeviceData(0), 63 sph_model->get_device_data(0), 64 current_time, 65 sph_particle_interaction_f, 66 udf_data, 67 handler); 68 } 69 else 70 { 71 SPHExplicitHandler<OMPBackend, false> handler { solver_model.GetDeviceData(0), 72 sph_model->get_device_data(0), 73 current_time, 74 it_type }; 75 device::launch_sph_interactions_particles< 76 OMPBackend, 77 Operator, 78 decltype(sph_particle_interaction_f)>( 79 solver_model.GetDeviceData(0), 80 sph_model->get_device_data(0), 81 current_time, 82 sph_particle_interaction_f, 83 udf_data, 84 handler); 85 } 86 } 87 else 88 { 89 if (solver_model.gen.nprt_active > 0) 90 { 91 if (solver_model.gen.periodic_domain.enabled()) 92 { 93 SPHApplyHandler<OMPBackend, true, true> handler { solver_model.GetDeviceData(0), 94 sph_model->get_device_data(0), 95 current_time, 96 it_type }; 97 device::launch_sph_interactions_particles< 98 OMPBackend, 99 Operator, 100 decltype(sph_particle_interaction_f)>( 101 solver_model.GetDeviceData(0), 102 sph_model->get_device_data(0), 103 current_time, 104 sph_particle_interaction_f, 105 udf_data, 106 handler); 107 } 108 else 109 { 110 SPHApplyHandler<OMPBackend, true, false> handler { 111 solver_model.GetDeviceData(0), 112 sph_model->get_device_data(0), 113 current_time, 114 it_type 115 }; 116 device::launch_sph_interactions_particles< 117 OMPBackend, 118 Operator, 119 decltype(sph_particle_interaction_f)>( 120 solver_model.GetDeviceData(0), 121 sph_model->get_device_data(0), 122 current_time, 123 sph_particle_interaction_f, 124 udf_data, 125 handler); 126 } 127 } 128 else 129 { 130 if (solver_model.gen.periodic_domain.enabled()) 131 { 132 SPHApplyHandler<OMPBackend, false, true> handler { 133 solver_model.GetDeviceData(0), 134 sph_model->get_device_data(0), 135 current_time, 136 it_type 137 }; 138 device::launch_sph_interactions_particles< 139 OMPBackend, 140 Operator, 141 decltype(sph_particle_interaction_f)>( 142 solver_model.GetDeviceData(0), 143 sph_model->get_device_data(0), 144 current_time, 145 sph_particle_interaction_f, 146 udf_data, 147 handler); 148 } 149 else 150 { 151 SPHApplyHandler<OMPBackend, false, false> handler { 152 solver_model.GetDeviceData(0), 153 sph_model->get_device_data(0), 154 current_time, 155 it_type 156 }; 157 device::launch_sph_interactions_particles< 158 OMPBackend, 159 Operator, 160 decltype(sph_particle_interaction_f)>( 161 solver_model.GetDeviceData(0), 162 sph_model->get_device_data(0), 163 current_time, 164 sph_particle_interaction_f, 165 udf_data, 166 handler); 167 } 168 } 169 } 170 } 171#endif 172} 173 174template <typename Operator = device::DefaultTrianglesOperator<OMPBackend>, typename Functor> 175inline void dispatch_sph_interactions_triangles_kernel( 176 IRockyModel& model, Functor& sph_triangle_interaction_f, void* udf_data) 177{ 178 SPHModel* sph_model = &(*model._impl->sph_model); 179 SPHDeviceModel* sph_device = model._impl->sph_model->get_device_data(0); 180 SDeviceModel* device_data = model.get_device_data(0); 181 182 if (sph_model->gen_sph.nprt_active <= 0) 183 return; 184 185#ifndef GPU_KERNEL_COMPILING_STAGE 186 #pragma omp parallel num_threads(model.get_number_of_processes()) 187 { 188 // SPH -> triangles 189 device::launch_sph_interactions_triangles<OMPBackend, Operator>( 190 device_data, 191 sph_device, 192 model.get_current_time(), 193 sph_triangle_interaction_f, 194 udf_data); 195 } 196#endif 197} 198 199} // namespace rocky::omp IRockyModelDefinition rocky_api.h:58 IRockyModel::get_current_timedouble get_current_time() constDefinition rocky_api.cpp:143