## On this page

  

 

 # api\_device\_joints

 Last update: 17.07.2025 

<a id="l00001" name="l00001"></a> 1\#pragma once

<a id="l00002" name="l00002"></a> 2 

<a id="l00003" name="l00003"></a> 3\#include &lt;rocky20/device/device\_model.hpp&gt;

<a id="l00004" name="l00004"></a> 4 

<a id="l00005" name="l00005"></a> 5namespace device

<a id="l00006" name="l00006"></a> 6{

<a id="l00007" name="l00007"></a> 7 

<a id="l00008" name="l00008"></a> 8template &lt;class backend, typename Callable&gt;

<a id="l00009" name="l00009"></a> 9__global__ void execute_joints_function(

<a id="l00010" name="l00010"></a> 10 SDeviceModel *device_data, int device_id, cu_real current_time,

<a id="l00011" name="l00011"></a> 11 Callable process_joint_ptr, void *udf_data

<a id="l00012" name="l00012"></a> 12)

<a id="l00013" name="l00013"></a> 13{

<a id="l00014" name="l00014"></a> 14 [IRockyDeviceModel](structIRockyDeviceModel.xhtml) device_model_api;

<a id="l00015" name="l00015"></a> 15 device_model_api._impl.model = device_data;

<a id="l00016" name="l00016"></a> 16 device_model_api._impl.current_time = current_time;

<a id="l00017" name="l00017"></a> 17 

<a id="l00018" name="l00018"></a> 18 [IRockyJoint](structIRockyJoint.xhtml) joint_api(device_data, device_id);

<a id="l00019" name="l00019"></a> 19 

<a id="l00020" name="l00020"></a> 20 // Thread

<a id="l00021" name="l00021"></a> 21 const cuda_general &amp;general = device_data-&gt;get_general();

<a id="l00022" name="l00022"></a> 22 int thread_start, thread_step, thread_stop;

<a id="l00023" name="l00023"></a> 23 backend::ThreadIteration(thread_start, thread_step, thread_stop, general.njn_prt_active);

<a id="l00024" name="l00024"></a> 24 for (int joint_index = thread_start; joint_index &lt; thread_stop; joint_index += thread_step)

<a id="l00025" name="l00025"></a> 25 {

<a id="l00026" name="l00026"></a> 26 joint_api._impl.reset(joint_index);

<a id="l00027" name="l00027"></a> 27 if (joint_api._impl.is_enabled())

<a id="l00028" name="l00028"></a> 28 {

<a id="l00029" name="l00029"></a> 29 //Process joint forces

<a id="l00030" name="l00030"></a> 30 process_joint_ptr(device_model_api, joint_api, udf_data);

<a id="l00031" name="l00031"></a> 31 }

<a id="l00032" name="l00032"></a> 32 }

<a id="l00033" name="l00033"></a> 33}

<a id="l00034" name="l00034"></a> 34 

<a id="l00035" name="l00035"></a> 35template &lt;class backend, typename Callable&gt;

<a id="l00036" name="l00036"></a> 36__host__ __device__ void compute_joint_forces(

<a id="l00037" name="l00037"></a> 37 SDeviceModel *device_data, int device_id,

<a id="l00038" name="l00038"></a> 38 Callable process_joint_ptr, void *udf_data

<a id="l00039" name="l00039"></a> 39)

<a id="l00040" name="l00040"></a> 40{

<a id="l00041" name="l00041"></a> 41 //General data

<a id="l00042" name="l00042"></a> 42 const cuda_general &amp;general = device_data-&gt;get_general();

<a id="l00043" name="l00043"></a> 43 const auto &amp;periodic = device_data-&gt;gen-&gt;periodic_domain;

<a id="l00044" name="l00044"></a> 44 bool is_cylindrical = periodic.is_cylindrical();

<a id="l00045" name="l00045"></a> 45 

<a id="l00046" name="l00046"></a> 46 [IRockyJoint](structIRockyJoint.xhtml) joint_api(device_data, device_id);

<a id="l00047" name="l00047"></a> 47 auto* data_recorder = device_data-&gt;joint_data_recorder;

<a id="l00048" name="l00048"></a> 48 

<a id="l00049" name="l00049"></a> 49 // Thread

<a id="l00050" name="l00050"></a> 50 int thread_start, thread_step, thread_stop;

<a id="l00051" name="l00051"></a> 51 backend::ThreadIteration(thread_start, thread_step, thread_stop, general.njn_prt_active);

<a id="l00052" name="l00052"></a> 52 for (int joint_index = thread_start; joint_index &lt; thread_stop; joint_index += thread_step)

<a id="l00053" name="l00053"></a> 53 {

<a id="l00054" name="l00054"></a> 54 if (joint_api._impl.get_joint_model(joint_index) == jmCustom)

<a id="l00055" name="l00055"></a> 55 {

<a id="l00056" name="l00056"></a> 56 joint_api._impl.reset(joint_index);

<a id="l00057" name="l00057"></a> 57 if (joint_api._impl.is_enabled())

<a id="l00058" name="l00058"></a> 58 {

<a id="l00059" name="l00059"></a> 59 //Process joint forces

<a id="l00060" name="l00060"></a> 60 

<a id="l00061" name="l00061"></a> 61 IRockyJointOutputData output_data{};

<a id="l00062" name="l00062"></a> 62 process_joint_ptr(joint_api, output_data, udf_data);

<a id="l00063" name="l00063"></a> 63 

<a id="l00064" name="l00064"></a> 64 auto &amp;joint = device_data-&gt;jnt_prt[joint_index];

<a id="l00065" name="l00065"></a> 65 

<a id="l00066" name="l00066"></a> 66 //Add joint forces and moments - home particle

<a id="l00067" name="l00067"></a> 67 cuda_particle_force &amp;h_particle_forces = device_data-&gt;frc[joint.nh];

<a id="l00068" name="l00068"></a> 68 

<a id="l00069" name="l00069"></a> 69 const auto&amp; force = output_data._force;

<a id="l00070" name="l00070"></a> 70 const auto&amp; home_moment = output_data._home_moment;

<a id="l00071" name="l00071"></a> 71 cu_real fxh = force.x;

<a id="l00072" name="l00072"></a> 72 cu_real fyh = force.y;

<a id="l00073" name="l00073"></a> 73 cu_real fzh = force.z;

<a id="l00074" name="l00074"></a> 74 cu_real mxh = home_moment.x;

<a id="l00075" name="l00075"></a> 75 cu_real myh = home_moment.y;

<a id="l00076" name="l00076"></a> 76 cu_real mzh = home_moment.z;

<a id="l00077" name="l00077"></a> 77 

<a id="l00078" name="l00078"></a> 78 if (is_cylindrical)

<a id="l00079" name="l00079"></a> 79 {

<a id="l00080" name="l00080"></a> 80 periodic.get_cylindric_periodic_position(

<a id="l00081" name="l00081"></a> 81 fxh, fyh, fzh, joint_api._impl.state.periodic_position, {}, {}, true);

<a id="l00082" name="l00082"></a> 82 periodic.get_cylindric_periodic_position(

<a id="l00083" name="l00083"></a> 83 mxh, myh, mzh, joint_api._impl.state.periodic_position, {}, {}, true);

<a id="l00084" name="l00084"></a> 84 }

<a id="l00085" name="l00085"></a> 85 

<a id="l00086" name="l00086"></a> 86 backend::AtomicAdd(&amp;h_particle_forces.fx, fxh);

<a id="l00087" name="l00087"></a> 87 backend::AtomicAdd(&amp;h_particle_forces.fy, fyh);

<a id="l00088" name="l00088"></a> 88 backend::AtomicAdd(&amp;h_particle_forces.fz, fzh);

<a id="l00089" name="l00089"></a> 89 backend::AtomicAdd(&amp;h_particle_forces.mx, mxh);

<a id="l00090" name="l00090"></a> 90 backend::AtomicAdd(&amp;h_particle_forces.my, myh);

<a id="l00091" name="l00091"></a> 91 backend::AtomicAdd(&amp;h_particle_forces.mz, mzh);

<a id="l00092" name="l00092"></a> 92 

<a id="l00093" name="l00093"></a> 93 //Add joint forces and moments - near particle

<a id="l00094" name="l00094"></a> 94 auto&amp; n_particle_force = device_data-&gt;frc[joint.nn];

<a id="l00095" name="l00095"></a> 95 

<a id="l00096" name="l00096"></a> 96 backend::AtomicAdd(&amp;(n_particle_force.fx), -force.x);

<a id="l00097" name="l00097"></a> 97 backend::AtomicAdd(&amp;(n_particle_force.fy), -force.y);

<a id="l00098" name="l00098"></a> 98 backend::AtomicAdd(&amp;(n_particle_force.fz), -force.z);

<a id="l00099" name="l00099"></a> 99 

<a id="l00100" name="l00100"></a> 100 const auto &amp;near_moment = output_data._near_moment;

<a id="l00101" name="l00101"></a> 101 backend::AtomicAdd(&amp;(n_particle_force.mx), near_moment.x);

<a id="l00102" name="l00102"></a> 102 backend::AtomicAdd(&amp;(n_particle_force.my), near_moment.y);

<a id="l00103" name="l00103"></a> 103 backend::AtomicAdd(&amp;(n_particle_force.mz), near_moment.z);

<a id="l00104" name="l00104"></a> 104 

<a id="l00105" name="l00105"></a> 105 if (data_recorder &amp;&amp; data_recorder-&gt;enabled)

<a id="l00106" name="l00106"></a> 106 {

<a id="l00107" name="l00107"></a> 107 const cu_real dt = device_data-&gt;gen-&gt;dt;

<a id="l00108" name="l00108"></a> 108 

<a id="l00109" name="l00109"></a> 109 // register forces

<a id="l00110" name="l00110"></a> 110 const cu_real3&amp; normal = joint_api._impl.state.normal_vector;

<a id="l00111" name="l00111"></a> 111 const cu_real normal_force = rocky::vector::dot(force, normal);

<a id="l00112" name="l00112"></a> 112 const cu_real tangential_force = rocky::vector::get_norm(force - normal_force * normal);

<a id="l00113" name="l00113"></a> 113 data_recorder-&gt;register_joint_forces(

<a id="l00114" name="l00114"></a> 114 device_data-&gt;joint_scalars,

<a id="l00115" name="l00115"></a> 115 normal_force, tangential_force,

<a id="l00116" name="l00116"></a> 116 joint_index, dt

<a id="l00117" name="l00117"></a> 117 );

<a id="l00118" name="l00118"></a> 118 

<a id="l00119" name="l00119"></a> 119 // register moments

<a id="l00120" name="l00120"></a> 120 const cu_real3 avg_moment = 0.5 * (home_moment + near_moment);

<a id="l00121" name="l00121"></a> 121 for (int i = 0; i &lt; 3; i++)

<a id="l00122" name="l00122"></a> 122 {

<a id="l00123" name="l00123"></a> 123 const cu_real3&amp; axis = joint_api._impl.state.ax[i];

<a id="l00124" name="l00124"></a> 124 const cu_real m_component = rocky::vector::dot(avg_moment, axis);

<a id="l00125" name="l00125"></a> 125 data_recorder-&gt;register_joint_moment(

<a id="l00126" name="l00126"></a> 126 device_data-&gt;joint_scalars,

<a id="l00127" name="l00127"></a> 127 static\_cast&lt;JointDataRecorder::JointMomentType&gt;(i),

<a id="l00128" name="l00128"></a> 128 m_component,

<a id="l00129" name="l00129"></a> 129 joint_index, dt

<a id="l00130" name="l00130"></a> 130 );

<a id="l00131" name="l00131"></a> 131 }

<a id="l00132" name="l00132"></a> 132 

<a id="l00133" name="l00133"></a> 133 // register stresses

<a id="l00134" name="l00134"></a> 134 if (data_recorder-&gt;stresses_enabled)

<a id="l00135" name="l00135"></a> 135 {

<a id="l00136" name="l00136"></a> 136 const cu_real area = joint_api._impl.state.area;

<a id="l00137" name="l00137"></a> 137 const cu_real tensile_stress = rocky::dmax(0.0, normal_force) / area;

<a id="l00138" name="l00138"></a> 138 const cu_real shear_stress = tangential_force / area;

<a id="l00139" name="l00139"></a> 139 const cu_real von_mises_stress = sqrt(tensile_stress * tensile_stress + 3.0 * shear_stress * shear_stress);

<a id="l00140" name="l00140"></a> 140 data_recorder-&gt;register_joint_stresses(

<a id="l00141" name="l00141"></a> 141 device_data-&gt;joint_scalars,

<a id="l00142" name="l00142"></a> 142 tensile_stress, shear_stress, von_mises_stress,

<a id="l00143" name="l00143"></a> 143 joint_index, dt

<a id="l00144" name="l00144"></a> 144 );

<a id="l00145" name="l00145"></a> 145 }

<a id="l00146" name="l00146"></a> 146 }

<a id="l00147" name="l00147"></a> 147 }

<a id="l00148" name="l00148"></a> 148 }

<a id="l00149" name="l00149"></a> 149 }

<a id="l00150" name="l00150"></a> 150}

<a id="l00151" name="l00151"></a> 151 

<a id="l00152" name="l00152"></a> 152template &lt;class backend, typename Callable&gt;

<a id="l00153" name="l00153"></a> 153__host__ __device__ void compute_joint_breakage(

<a id="l00154" name="l00154"></a> 154 SDeviceModel* device_data, int device_id,

<a id="l00155" name="l00155"></a> 155 Callable process_joint_breakage_ptr, void* udf_data,

<a id="l00156" name="l00156"></a> 156 typename backend::atomic_int_ptr nbreak_joints_a, cu_real time_curr

<a id="l00157" name="l00157"></a> 157)

<a id="l00158" name="l00158"></a> 158{

<a id="l00159" name="l00159"></a> 159 const cuda_general&amp; general = device_data-&gt;get_general();

<a id="l00160" name="l00160"></a> 160 

<a id="l00161" name="l00161"></a> 161 int njb_old = 0;

<a id="l00162" name="l00162"></a> 162 [IRockyJoint](structIRockyJoint.xhtml) joint_api(device_data, device_id);

<a id="l00163" name="l00163"></a> 163 auto&amp; rocky_joint = joint_api._impl;

<a id="l00164" name="l00164"></a> 164 

<a id="l00165" name="l00165"></a> 165 int thread_start, thread_step, thread_stop;

<a id="l00166" name="l00166"></a> 166 backend::ThreadIteration(thread_start, thread_step, thread_stop, general.njn_prt_brk_active);

<a id="l00167" name="l00167"></a> 167 for (int brjnt_index = thread_start; brjnt_index &lt; thread_stop; brjnt_index += thread_step)

<a id="l00168" name="l00168"></a> 168 {

<a id="l00169" name="l00169"></a> 169 int joint_index = device_data-&gt;breakable_joints[brjnt_index] - general.njn_prt_active_start;

<a id="l00170" name="l00170"></a> 170 

<a id="l00171" name="l00171"></a> 171 if (rocky_joint.get_breakage_model(joint_index) == pbmDiscreteCustom)

<a id="l00172" name="l00172"></a> 172 {

<a id="l00173" name="l00173"></a> 173 rocky_joint.reset(joint_index);

<a id="l00174" name="l00174"></a> 174 bool broken = false;

<a id="l00175" name="l00175"></a> 175 

<a id="l00176" name="l00176"></a> 176 if (rocky_joint.is_enabled())

<a id="l00177" name="l00177"></a> 177 {

<a id="l00178" name="l00178"></a> 178 if (time_curr - rocky_joint.get_particle_release_time() &lt; general.breakage_start_release_time)

<a id="l00179" name="l00179"></a> 179 return;

<a id="l00180" name="l00180"></a> 180 

<a id="l00181" name="l00181"></a> 181 process_joint_breakage_ptr(joint_api, broken, udf_data);

<a id="l00182" name="l00182"></a> 182 

<a id="l00183" name="l00183"></a> 183 if (broken)

<a id="l00184" name="l00184"></a> 184 {

<a id="l00185" name="l00185"></a> 185 backend::AtomicAdd(nbreak_joints_a, 1, njb_old);

<a id="l00186" name="l00186"></a> 186 device_data-&gt;broken_joints[njb_old] = joint_index;

<a id="l00187" name="l00187"></a> 187 device_data-&gt;joint_scalars-&gt;get_breakage_state(joint_index) = dbsFullyBroken;

<a id="l00188" name="l00188"></a> 188 }

<a id="l00189" name="l00189"></a> 189 }

<a id="l00190" name="l00190"></a> 190 }

<a id="l00191" name="l00191"></a> 191 }

<a id="l00192" name="l00192"></a> 192}

<a id="l00193" name="l00193"></a> 193 

<a id="l00194" name="l00194"></a> 194} // namespace device

[IRockyDeviceModel](structIRockyDeviceModel.xhtml)

**Definition** rocky_model_api.hpp:19



[IRockyJoint](structIRockyJoint.xhtml)

**Definition** rocky_joint_api.hpp:32