api_device_joints Last update: 17.07.2025 1#pragma once 2 3#include <rocky20/device/device_model.hpp> 4 5namespace device 6{ 7 8template <class backend, typename Callable> 9__global__ void execute_joints_function( 10 SDeviceModel *device_data, int device_id, cu_real current_time, 11 Callable process_joint_ptr, void *udf_data 12) 13{ 14 IRockyDeviceModel device_model_api; 15 device_model_api._impl.model = device_data; 16 device_model_api._impl.current_time = current_time; 17 18 IRockyJoint joint_api(device_data, device_id); 19 20 // Thread 21 const cuda_general &general = device_data->get_general(); 22 int thread_start, thread_step, thread_stop; 23 backend::ThreadIteration(thread_start, thread_step, thread_stop, general.njn_prt_active); 24 for (int joint_index = thread_start; joint_index < thread_stop; joint_index += thread_step) 25 { 26 joint_api._impl.reset(joint_index); 27 if (joint_api._impl.is_enabled()) 28 { 29 //Process joint forces 30 process_joint_ptr(device_model_api, joint_api, udf_data); 31 } 32 } 33} 34 35template <class backend, typename Callable> 36__host__ __device__ void compute_joint_forces( 37 SDeviceModel *device_data, int device_id, 38 Callable process_joint_ptr, void *udf_data 39) 40{ 41 //General data 42 const cuda_general &general = device_data->get_general(); 43 const auto &periodic = device_data->gen->periodic_domain; 44 bool is_cylindrical = periodic.is_cylindrical(); 45 46 IRockyJoint joint_api(device_data, device_id); 47 auto* data_recorder = device_data->joint_data_recorder; 48 49 // Thread 50 int thread_start, thread_step, thread_stop; 51 backend::ThreadIteration(thread_start, thread_step, thread_stop, general.njn_prt_active); 52 for (int joint_index = thread_start; joint_index < thread_stop; joint_index += thread_step) 53 { 54 if (joint_api._impl.get_joint_model(joint_index) == jmCustom) 55 { 56 joint_api._impl.reset(joint_index); 57 if (joint_api._impl.is_enabled()) 58 { 59 //Process joint forces 60 61 IRockyJointOutputData output_data{}; 62 process_joint_ptr(joint_api, output_data, udf_data); 63 64 auto &joint = device_data->jnt_prt[joint_index]; 65 66 //Add joint forces and moments - home particle 67 cuda_particle_force &h_particle_forces = device_data->frc[joint.nh]; 68 69 const auto& force = output_data._force; 70 const auto& home_moment = output_data._home_moment; 71 cu_real fxh = force.x; 72 cu_real fyh = force.y; 73 cu_real fzh = force.z; 74 cu_real mxh = home_moment.x; 75 cu_real myh = home_moment.y; 76 cu_real mzh = home_moment.z; 77 78 if (is_cylindrical) 79 { 80 periodic.get_cylindric_periodic_position( 81 fxh, fyh, fzh, joint_api._impl.state.periodic_position, {}, {}, true); 82 periodic.get_cylindric_periodic_position( 83 mxh, myh, mzh, joint_api._impl.state.periodic_position, {}, {}, true); 84 } 85 86 backend::AtomicAdd(&h_particle_forces.fx, fxh); 87 backend::AtomicAdd(&h_particle_forces.fy, fyh); 88 backend::AtomicAdd(&h_particle_forces.fz, fzh); 89 backend::AtomicAdd(&h_particle_forces.mx, mxh); 90 backend::AtomicAdd(&h_particle_forces.my, myh); 91 backend::AtomicAdd(&h_particle_forces.mz, mzh); 92 93 //Add joint forces and moments - near particle 94 auto& n_particle_force = device_data->frc[joint.nn]; 95 96 backend::AtomicAdd(&(n_particle_force.fx), -force.x); 97 backend::AtomicAdd(&(n_particle_force.fy), -force.y); 98 backend::AtomicAdd(&(n_particle_force.fz), -force.z); 99 100 const auto &near_moment = output_data._near_moment; 101 backend::AtomicAdd(&(n_particle_force.mx), near_moment.x); 102 backend::AtomicAdd(&(n_particle_force.my), near_moment.y); 103 backend::AtomicAdd(&(n_particle_force.mz), near_moment.z); 104 105 if (data_recorder && data_recorder->enabled) 106 { 107 const cu_real dt = device_data->gen->dt; 108 109 // register forces 110 const cu_real3& normal = joint_api._impl.state.normal_vector; 111 const cu_real normal_force = rocky::vector::dot(force, normal); 112 const cu_real tangential_force = rocky::vector::get_norm(force - normal_force * normal); 113 data_recorder->register_joint_forces( 114 device_data->joint_scalars, 115 normal_force, tangential_force, 116 joint_index, dt 117 ); 118 119 // register moments 120 const cu_real3 avg_moment = 0.5 * (home_moment + near_moment); 121 for (int i = 0; i < 3; i++) 122 { 123 const cu_real3& axis = joint_api._impl.state.ax[i]; 124 const cu_real m_component = rocky::vector::dot(avg_moment, axis); 125 data_recorder->register_joint_moment( 126 device_data->joint_scalars, 127 static_cast<JointDataRecorder::JointMomentType>(i), 128 m_component, 129 joint_index, dt 130 ); 131 } 132 133 // register stresses 134 if (data_recorder->stresses_enabled) 135 { 136 const cu_real area = joint_api._impl.state.area; 137 const cu_real tensile_stress = rocky::dmax(0.0, normal_force) / area; 138 const cu_real shear_stress = tangential_force / area; 139 const cu_real von_mises_stress = sqrt(tensile_stress * tensile_stress + 3.0 * shear_stress * shear_stress); 140 data_recorder->register_joint_stresses( 141 device_data->joint_scalars, 142 tensile_stress, shear_stress, von_mises_stress, 143 joint_index, dt 144 ); 145 } 146 } 147 } 148 } 149 } 150} 151 152template <class backend, typename Callable> 153__host__ __device__ void compute_joint_breakage( 154 SDeviceModel* device_data, int device_id, 155 Callable process_joint_breakage_ptr, void* udf_data, 156 typename backend::atomic_int_ptr nbreak_joints_a, cu_real time_curr 157) 158{ 159 const cuda_general& general = device_data->get_general(); 160 161 int njb_old = 0; 162 IRockyJoint joint_api(device_data, device_id); 163 auto& rocky_joint = joint_api._impl; 164 165 int thread_start, thread_step, thread_stop; 166 backend::ThreadIteration(thread_start, thread_step, thread_stop, general.njn_prt_brk_active); 167 for (int brjnt_index = thread_start; brjnt_index < thread_stop; brjnt_index += thread_step) 168 { 169 int joint_index = device_data->breakable_joints[brjnt_index] - general.njn_prt_active_start; 170 171 if (rocky_joint.get_breakage_model(joint_index) == pbmDiscreteCustom) 172 { 173 rocky_joint.reset(joint_index); 174 bool broken = false; 175 176 if (rocky_joint.is_enabled()) 177 { 178 if (time_curr - rocky_joint.get_particle_release_time() < general.breakage_start_release_time) 179 return; 180 181 process_joint_breakage_ptr(joint_api, broken, udf_data); 182 183 if (broken) 184 { 185 backend::AtomicAdd(nbreak_joints_a, 1, njb_old); 186 device_data->broken_joints[njb_old] = joint_index; 187 device_data->joint_scalars->get_breakage_state(joint_index) = dbsFullyBroken; 188 } 189 } 190 } 191 } 192} 193 194} // namespace device IRockyDeviceModelDefinition rocky_model_api.hpp:19 IRockyJointDefinition rocky_joint_api.hpp:32
