## On this page

  

 

 # api\_device\_sph

 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/api/rocky\_api\_sph\_functions.h&gt;

<a id="l00004" name="l00004"></a> 4\#include &lt;rocky20/api/rocky\_model\_api.hpp&gt;

<a id="l00005" name="l00005"></a> 5\#include &lt;rocky20/api/sph/rocky\_sph\_element\_api.hpp&gt;

<a id="l00006" name="l00006"></a> 6\#include &lt;rocky20/api/sph/rocky\_sph\_interaction\_api.hpp&gt;

<a id="l00007" name="l00007"></a> 7\#include &lt;rocky20/api/sph/rocky\_sph\_model\_api.hpp&gt;

<a id="l00008" name="l00008"></a> 8 

<a id="l00009" name="l00009"></a> 9namespace device {

<a id="l00010" name="l00010"></a> 10 

<a id="l00011" name="l00011"></a> 11template &lt;class Backend&gt;

<a id="l00012" name="l00012"></a> 12struct DefaultElementsOperator

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

<a id="l00014" name="l00014"></a> 14 template &lt;typename Callable&gt;

<a id="l00015" name="l00015"></a> 15 __host__ __device__ void operator()(

<a id="l00016" name="l00016"></a> 16 Callable function_ptr,

<a id="l00017" name="l00017"></a> 17 [IRockyDeviceModel](structIRockyDeviceModel.xhtml)&amp; device_model_api,

<a id="l00018" name="l00018"></a> 18 [IRockySPHDeviceModel](structIRockySPHDeviceModel.xhtml)&amp; sph_model,

<a id="l00019" name="l00019"></a> 19 [IRockySPHElementInteraction](structIRockySPHElementInteraction.xhtml)&amp; sph_interaction_api,

<a id="l00020" name="l00020"></a> 20 void* udf_data)

<a id="l00021" name="l00021"></a> 21 {

<a id="l00022" name="l00022"></a> 22 function_ptr(device_model_api, sph_model, sph_interaction_api, udf_data);

<a id="l00023" name="l00023"></a> 23 }

<a id="l00024" name="l00024"></a> 24};

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

<a id="l00026" name="l00026"></a> 26template &lt;class Backend&gt;

<a id="l00027" name="l00027"></a> 27struct DefaultTrianglesOperator

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

<a id="l00029" name="l00029"></a> 29 template &lt;typename Callable&gt;

<a id="l00030" name="l00030"></a> 30 __host__ __device__ void operator()(

<a id="l00031" name="l00031"></a> 31 Callable function_ptr,

<a id="l00032" name="l00032"></a> 32 [IRockyDeviceModel](structIRockyDeviceModel.xhtml)&amp; device_model_api,

<a id="l00033" name="l00033"></a> 33 [IRockySPHDeviceModel](structIRockySPHDeviceModel.xhtml)&amp; sph_model,

<a id="l00034" name="l00034"></a> 34 [IRockySPHTriangleInteraction](structIRockySPHTriangleInteraction.xhtml)&amp; sph_interaction_api,

<a id="l00035" name="l00035"></a> 35 void* udf_data)

<a id="l00036" name="l00036"></a> 36 {

<a id="l00037" name="l00037"></a> 37 function_ptr(device_model_api, sph_model, sph_interaction_api, udf_data);

<a id="l00038" name="l00038"></a> 38 }

<a id="l00039" name="l00039"></a> 39};

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

<a id="l00041" name="l00041"></a> 41template &lt;class Backend&gt;

<a id="l00042" name="l00042"></a> 42struct DefaultOperator

<a id="l00043" name="l00043"></a> 43{

<a id="l00044" name="l00044"></a> 44 template &lt;typename Callable&gt;

<a id="l00045" name="l00045"></a> 45 __host__ __device__ void operator()(

<a id="l00046" name="l00046"></a> 46 Callable function_ptr,

<a id="l00047" name="l00047"></a> 47 [IRockyDeviceModel](structIRockyDeviceModel.xhtml)&amp; device_model_api,

<a id="l00048" name="l00048"></a> 48 [IRockySPHDeviceModel](structIRockySPHDeviceModel.xhtml)&amp; sph_model,

<a id="l00049" name="l00049"></a> 49 [IRockySPHElement](structIRockySPHElement.xhtml)&amp; sph_element_api,

<a id="l00050" name="l00050"></a> 50 void* udf_data)

<a id="l00051" name="l00051"></a> 51 {

<a id="l00052" name="l00052"></a> 52 function_ptr(device_model_api, sph_model, sph_element_api, udf_data);

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

<a id="l00054" name="l00054"></a> 54};

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

<a id="l00056" name="l00056"></a> 56template &lt;class Backend&gt;

<a id="l00057" name="l00057"></a> 57struct HeatTransferElementsOperator

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

<a id="l00059" name="l00059"></a> 59 template &lt;typename Callable&gt;

<a id="l00060" name="l00060"></a> 60 __host__ __device__ void operator()(

<a id="l00061" name="l00061"></a> 61 Callable function_ptr,

<a id="l00062" name="l00062"></a> 62 [IRockyDeviceModel](structIRockyDeviceModel.xhtml)&amp; device_model_api,

<a id="l00063" name="l00063"></a> 63 [IRockySPHDeviceModel](structIRockySPHDeviceModel.xhtml)&amp; sph_model,

<a id="l00064" name="l00064"></a> 64 [IRockySPHElementInteraction](structIRockySPHElementInteraction.xhtml)&amp; sph_interaction_api,

<a id="l00065" name="l00065"></a> 65 void* udf_data)

<a id="l00066" name="l00066"></a> 66 {

<a id="l00067" name="l00067"></a> 67 // Compute and add heat transfer

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

<a id="l00069" name="l00069"></a> 69 function_ptr(device_model_api, sph_model, sph_interaction_api, udf_data);

<a id="l00070" name="l00070"></a> 70 auto home_sph_element = sph_interaction_api.[get\_home\_element](structIRockySPHElementInteraction.xhtml#a7eb0236ba9b25b58b0ca258b1e79976b)();

<a id="l00071" name="l00071"></a> 71 auto near_sph_element = sph_interaction_api.[get\_near\_element](structIRockySPHElementInteraction.xhtml#a0e9d2687e1c8a66357730003f850fe26)();

<a id="l00072" name="l00072"></a> 72 home_sph_element.[get\_scalars](structIRockySPHElement.xhtml#a326bd60795936f9c4b77a9e000b1fa99)().[add\_heat\_transfer](structIRockySPHElementScalars.xhtml#ace6ee0291cca05f39eb04ba4c760afe0)(heat_transfer);

<a id="l00073" name="l00073"></a> 73 near_sph_element.get_scalars().add_heat_transfer(-heat_transfer);

<a id="l00074" name="l00074"></a> 74 

<a id="l00075" name="l00075"></a> 75 // Recording coupled heat transfer

<a id="l00076" name="l00076"></a> 76 if (home_sph_element.is_dem_coupled())

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

<a id="l00078" name="l00078"></a> 78 auto dem_particle = home_sph_element.get_linked_dem_particle();

<a id="l00079" name="l00079"></a> 79 auto dem_scalars = dem_particle.get_transfer_scalars();

<a id="l00080" name="l00080"></a> 80 dem_scalars.add_heat_transfer(heat_transfer);

<a id="l00081" name="l00081"></a> 81 }

<a id="l00082" name="l00082"></a> 82 

<a id="l00083" name="l00083"></a> 83 if (near_sph_element.is_dem_coupled())

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

<a id="l00085" name="l00085"></a> 85 auto dem_particle = near_sph_element.get_linked_dem_particle();

<a id="l00086" name="l00086"></a> 86 auto dem_scalars = dem_particle.get_transfer_scalars();

<a id="l00087" name="l00087"></a> 87 dem_scalars.add_heat_transfer(-heat_transfer);

<a id="l00088" name="l00088"></a> 88 }

<a id="l00089" name="l00089"></a> 89 }

<a id="l00090" name="l00090"></a> 90};

<a id="l00091" name="l00091"></a> 91 

<a id="l00092" name="l00092"></a> 92template &lt;class Backend&gt;

<a id="l00093" name="l00093"></a> 93struct HeatTransferTrianglesOperator

<a id="l00094" name="l00094"></a> 94{

<a id="l00095" name="l00095"></a> 95 template &lt;typename Callable&gt;

<a id="l00096" name="l00096"></a> 96 __host__ __device__ void operator()(

<a id="l00097" name="l00097"></a> 97 Callable function_ptr,

<a id="l00098" name="l00098"></a> 98 [IRockyDeviceModel](structIRockyDeviceModel.xhtml)&amp; device_model_api,

<a id="l00099" name="l00099"></a> 99 [IRockySPHDeviceModel](structIRockySPHDeviceModel.xhtml)&amp; sph_model,

<a id="l00100" name="l00100"></a> 100 [IRockySPHTriangleInteraction](structIRockySPHTriangleInteraction.xhtml)&amp; sph_interaction_api,

<a id="l00101" name="l00101"></a> 101 void* udf_data)

<a id="l00102" name="l00102"></a> 102 {

<a id="l00103" name="l00103"></a> 103 // Compute and add heat transfer

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

<a id="l00105" name="l00105"></a> 105 function_ptr(device_model_api, sph_model, sph_interaction_api, udf_data);

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

<a id="l00107" name="l00107"></a> 107 // reduction factor for SPH elements with multiple triangle interactions

<a id="l00108" name="l00108"></a> 108 // (1 / number of simultaneous triangle interactions)

<a id="l00109" name="l00109"></a> 109 heat_transfer /= max(1.0f, sph_model._impl.device_model-&gt;bnd_force_coefficient[sph_interaction_api._impl-&gt;home_index]);

<a id="l00110" name="l00110"></a> 110 

<a id="l00111" name="l00111"></a> 111 sph_interaction_api.[get\_home\_element](structIRockySPHTriangleInteraction.xhtml#ae83d8750e222b2a46e6f901fd20b80d5)().[get\_scalars](structIRockySPHElement.xhtml#a326bd60795936f9c4b77a9e000b1fa99)().[add\_heat\_transfer](structIRockySPHElementScalars.xhtml#ace6ee0291cca05f39eb04ba4c760afe0)(heat_transfer);

<a id="l00112" name="l00112"></a> 112 

<a id="l00113" name="l00113"></a> 113 // Recording data for SPH Boundary Statistics

<a id="l00114" name="l00114"></a> 114 const auto&amp; data_recorder = sph_model._impl.device_model-&gt;boundary_data_recorder;

<a id="l00115" name="l00115"></a> 115 if (data_recorder

<a id="l00116" name="l00116"></a> 116 &amp;&amp; data_recorder-&gt;is_enabled(SPHBoundaryDataRecorder::tSPHTriangleHeatTransfer))

<a id="l00117" name="l00117"></a> 117 {

<a id="l00118" name="l00118"></a> 118 const int index =

<a id="l00119" name="l00119"></a> 119 data_recorder-&gt;get_scalar_index(SPHBoundaryDataRecorder::tSPHTriangleHeatTransfer);

<a id="l00120" name="l00120"></a> 120 auto triangle_scalars = sph_interaction_api.[get\_near\_triangle](structIRockySPHTriangleInteraction.xhtml#a469b4ac53234c9e707fee62725b0470a)().[get\_scalars](structIRockyGeometryTriangle.xhtml#a525086257224ba19ec3bbfec2522bdba)();

<a id="l00121" name="l00121"></a> 121 triangle_scalars.[add\_scalar](structIRockyTriangleScalars.xhtml#ab9721444212bd5fe4e3af03fe3a70294)(index, heat_transfer * device_model_api.[get\_timestep](structIRockyDeviceModel.xhtml#ab03adbadfc1d2a7a8f3a1f849be73d1a)());

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

<a id="l00123" name="l00123"></a> 123 

<a id="l00124" name="l00124"></a> 124 // Accumulation in a scalar used for heat coupling (positive when triangles receive heat)

<a id="l00125" name="l00125"></a> 125 auto* triangle_scalars = device_model_api._impl.model-&gt;triangle_scalars;

<a id="l00126" name="l00126"></a> 126 if (triangle_scalars &amp;&amp; triangle_scalars-&gt;is_heat_exchanged_enabled())

<a id="l00127" name="l00127"></a> 127 {

<a id="l00128" name="l00128"></a> 128 Backend::AtomicAdd(

<a id="l00129" name="l00129"></a> 129 &amp;triangle_scalars-&gt;get_heat_exchanged(sph_interaction_api._impl-&gt;near_index),

<a id="l00130" name="l00130"></a> 130 static\_cast&lt;double&gt;(-heat_transfer * device_model_api.[get\_timestep](structIRockyDeviceModel.xhtml#ab03adbadfc1d2a7a8f3a1f849be73d1a)()));

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

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

<a id="l00133" name="l00133"></a> 133};

<a id="l00134" name="l00134"></a> 134 

<a id="l00135" name="l00135"></a> 135template &lt;class Backend&gt;

<a id="l00136" name="l00136"></a> 136struct ForceElementInteractionOperator

<a id="l00137" name="l00137"></a> 137{

<a id="l00138" name="l00138"></a> 138 template &lt;typename Callable&gt;

<a id="l00139" name="l00139"></a> 139 __host__ __device__ void operator()(

<a id="l00140" name="l00140"></a> 140 Callable function_ptr,

<a id="l00141" name="l00141"></a> 141 [IRockyDeviceModel](structIRockyDeviceModel.xhtml)&amp; device_model_api,

<a id="l00142" name="l00142"></a> 142 [IRockySPHDeviceModel](structIRockySPHDeviceModel.xhtml)&amp; sph_model,

<a id="l00143" name="l00143"></a> 143 [IRockySPHElementInteraction](structIRockySPHElementInteraction.xhtml)&amp; sph_interaction_api,

<a id="l00144" name="l00144"></a> 144 void* udf_data)

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

<a id="l00146" name="l00146"></a> 146 // custom calculation of the element-to-element acceleration

<a id="l00147" name="l00147"></a> 147 const real_sph3 acceleration =

<a id="l00148" name="l00148"></a> 148 function_ptr(device_model_api, sph_model, sph_interaction_api, udf_data);

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

<a id="l00150" name="l00150"></a> 150 sph_interaction_api.[add\_acceleration](structIRockySPHElementInteraction.xhtml#a713381036ee74bb8476435037c14124b)(acceleration);

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

<a id="l00152" name="l00152"></a> 152};

<a id="l00153" name="l00153"></a> 153 

<a id="l00154" name="l00154"></a> 154template &lt;class Backend&gt;

<a id="l00155" name="l00155"></a> 155struct ForceTriangleInteractionOperator

<a id="l00156" name="l00156"></a> 156{

<a id="l00157" name="l00157"></a> 157 template &lt;typename Callable&gt;

<a id="l00158" name="l00158"></a> 158 __host__ __device__ void operator()(

<a id="l00159" name="l00159"></a> 159 Callable function_ptr,

<a id="l00160" name="l00160"></a> 160 [IRockyDeviceModel](structIRockyDeviceModel.xhtml)&amp; device_model_api,

<a id="l00161" name="l00161"></a> 161 [IRockySPHDeviceModel](structIRockySPHDeviceModel.xhtml)&amp; sph_model,

<a id="l00162" name="l00162"></a> 162 [IRockySPHTriangleInteraction](structIRockySPHTriangleInteraction.xhtml)&amp; sph_interaction_api,

<a id="l00163" name="l00163"></a> 163 void* udf_data)

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

<a id="l00165" name="l00165"></a> 165 // custom calculation of the element-to-triangle acceleration

<a id="l00166" name="l00166"></a> 166 const real_sph3 acceleration =

<a id="l00167" name="l00167"></a> 167 function_ptr(device_model_api, sph_model, sph_interaction_api, udf_data);

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

<a id="l00169" name="l00169"></a> 169 sph_interaction_api.[add\_acceleration](structIRockySPHTriangleInteraction.xhtml#a1578dfd6da083096eb169fe24913dd5b)(acceleration);

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

<a id="l00171" name="l00171"></a> 171};

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

<a id="l00173" name="l00173"></a> 173template &lt;class Backend&gt;

<a id="l00174" name="l00174"></a> 174struct ForceElementOperator

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

<a id="l00176" name="l00176"></a> 176 template &lt;typename Callable&gt;

<a id="l00177" name="l00177"></a> 177 __host__ __device__ void operator()(

<a id="l00178" name="l00178"></a> 178 Callable function_ptr,

<a id="l00179" name="l00179"></a> 179 [IRockyDeviceModel](structIRockyDeviceModel.xhtml)&amp; device_model_api,

<a id="l00180" name="l00180"></a> 180 [IRockySPHDeviceModel](structIRockySPHDeviceModel.xhtml)&amp; sph_model,

<a id="l00181" name="l00181"></a> 181 [IRockySPHElement](structIRockySPHElement.xhtml)&amp; sph_element_api,

<a id="l00182" name="l00182"></a> 182 void* udf_data)

<a id="l00183" name="l00183"></a> 183 {

<a id="l00184" name="l00184"></a> 184 // custom calculation of the element acceleration (due to a body force)

<a id="l00185" name="l00185"></a> 185 const real_sph3 acceleration =

<a id="l00186" name="l00186"></a> 186 function_ptr(device_model_api, sph_model, sph_element_api, udf_data);

<a id="l00187" name="l00187"></a> 187 

<a id="l00188" name="l00188"></a> 188 sph_element_api.[add\_acceleration](structIRockySPHElement.xhtml#a8f8e5c79577f1c799383d506d2415c2b)(acceleration);

<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> 192template &lt;class backend, typename Operator, typename Callable&gt;

<a id="l00193" name="l00193"></a> 193__host__ __device__ void launch_sph_particles_function(

<a id="l00194" name="l00194"></a> 194 SDeviceModel* device_data,

<a id="l00195" name="l00195"></a> 195 SPHDeviceModel* sph_device,

<a id="l00196" name="l00196"></a> 196 double current_time,

<a id="l00197" name="l00197"></a> 197 Callable function_ptr,

<a id="l00198" name="l00198"></a> 198 void* udf_data)

<a id="l00199" name="l00199"></a> 199{

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

<a id="l00201" name="l00201"></a> 201 device_model_api._impl.model = device_data;

<a id="l00202" name="l00202"></a> 202 device_model_api._impl.current_time = current_time;

<a id="l00203" name="l00203"></a> 203 [IRockySPHDeviceModel](structIRockySPHDeviceModel.xhtml) sph_model(device_data, sph_device-&gt;gen_sph, sph_device);

<a id="l00204" name="l00204"></a> 204 cuda_particle_sph* sph_elements = sph_device-&gt;prt_sph;

<a id="l00205" name="l00205"></a> 205 

<a id="l00206" name="l00206"></a> 206 int thread_start, thread_step, thread_stop;

<a id="l00207" name="l00207"></a> 207 backend::ThreadIteration(

<a id="l00208" name="l00208"></a> 208 thread_start, thread_step, thread_stop, sph_device-&gt;gen_sph-&gt;nprt_active);

<a id="l00209" name="l00209"></a> 209 for (int i = thread_start; i &lt; thread_stop; i += thread_step)

<a id="l00210" name="l00210"></a> 210 {

<a id="l00211" name="l00211"></a> 211 if (sph_elements[i].enabled() &amp;&amp; !sph_elements[i].ghost())

<a id="l00212" name="l00212"></a> 212 {

<a id="l00213" name="l00213"></a> 213 // Particle Interface

<a id="l00214" name="l00214"></a> 214 [IRockySPHElement](structIRockySPHElement.xhtml) sph_particle_api { device_data, sph_device, i };

<a id="l00215" name="l00215"></a> 215 

<a id="l00216" name="l00216"></a> 216 // Trigger user kernel

<a id="l00217" name="l00217"></a> 217 Operator caller{};

<a id="l00218" name="l00218"></a> 218 caller(function_ptr, device_model_api, sph_model, sph_particle_api, udf_data);

<a id="l00219" name="l00219"></a> 219 }

<a id="l00220" name="l00220"></a> 220 }

<a id="l00221" name="l00221"></a> 221}

<a id="l00222" name="l00222"></a> 222 

<a id="l00223" name="l00223"></a> 223template &lt;class Backend, typename Operator, typename Callable&gt;

<a id="l00224" name="l00224"></a> 224__host__ __device__ void launch_sph_interactions_particles(

<a id="l00225" name="l00225"></a> 225 SDeviceModel* device_data,

<a id="l00226" name="l00226"></a> 226 SPHDeviceModel* sph_device,

<a id="l00227" name="l00227"></a> 227 double current_time,

<a id="l00228" name="l00228"></a> 228 Callable function_ptr,

<a id="l00229" name="l00229"></a> 229 void* udf_data,

<a id="l00230" name="l00230"></a> 230 IterationType it_type)

<a id="l00231" name="l00231"></a> 231{

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

<a id="l00233" name="l00233"></a> 233 device_model_api._impl.model = device_data;

<a id="l00234" name="l00234"></a> 234 device_model_api._impl.current_time = current_time;

<a id="l00235" name="l00235"></a> 235 [IRockySPHDeviceModel](structIRockySPHDeviceModel.xhtml) sph_model(device_data, sph_device-&gt;gen_sph, sph_device);

<a id="l00236" name="l00236"></a> 236 

<a id="l00237" name="l00237"></a> 237 RockySPHElementInteraction sph_interaction(device_data, sph_device, current_time);

<a id="l00238" name="l00238"></a> 238 [IRockySPHElementInteraction](structIRockySPHElementInteraction.xhtml) sph_interaction_api { &amp;sph_interaction };

<a id="l00239" name="l00239"></a> 239 

<a id="l00240" name="l00240"></a> 240 cuda_general_sph* gen_sph = sph_device-&gt;gen_sph;

<a id="l00241" name="l00241"></a> 241 const auto&amp; periodic_domain = device_data-&gt;gen-&gt;periodic_domain;

<a id="l00242" name="l00242"></a> 242 

<a id="l00243" name="l00243"></a> 243 // Local storage for interactions and periodics

<a id="l00244" name="l00244"></a> 244 int* nn_list = nullptr;

<a id="l00245" name="l00245"></a> 245 PeriodicPosition* nn_periodic = nullptr;

<a id="l00246" name="l00246"></a> 246 if (!gen_sph-&gt;use_particles_neighbors_list)

<a id="l00247" name="l00247"></a> 247 {

<a id="l00248" name="l00248"></a> 248 nn_list = sph_device-&gt;local_n_list + Backend::ThreadIndex() * gen_sph-&gt;num_neighbors_max;

<a id="l00249" name="l00249"></a> 249 if (periodic_domain.enabled())

<a id="l00250" name="l00250"></a> 250 nn_periodic =

<a id="l00251" name="l00251"></a> 251 sph_device-&gt;local_n_periodic + Backend::ThreadIndex() * gen_sph-&gt;num_neighbors_max;

<a id="l00252" name="l00252"></a> 252 }

<a id="l00253" name="l00253"></a> 253 

<a id="l00254" name="l00254"></a> 254 // Thread

<a id="l00255" name="l00255"></a> 255 // TODO (ROCKY-20361): This will no longer be needed once we use SPH Search Handlers

<a id="l00256" name="l00256"></a> 256 const int n_entities = [&amp;gen_sph, it_type] {

<a id="l00257" name="l00257"></a> 257 if (!gen_sph-&gt;use_particles_neighbors_list)

<a id="l00258" name="l00258"></a> 258 return gen_sph-&gt;nprt_active;

<a id="l00259" name="l00259"></a> 259 if (it_type == itFreeAndFreeCoupled)

<a id="l00260" name="l00260"></a> 260 {

<a id="l00261" name="l00261"></a> 261 return gen_sph-&gt;ncont_ff_p_int + gen_sph-&gt;ncont_fc_p_int;

<a id="l00262" name="l00262"></a> 262 }

<a id="l00263" name="l00263"></a> 263 else if (it_type == itFreeCoupledOnly)

<a id="l00264" name="l00264"></a> 264 {

<a id="l00265" name="l00265"></a> 265 return gen_sph-&gt;ncont_fc_p_int;

<a id="l00266" name="l00266"></a> 266 }

<a id="l00267" name="l00267"></a> 267 else // itAll

<a id="l00268" name="l00268"></a> 268 {

<a id="l00269" name="l00269"></a> 269 return gen_sph-&gt;ncont_ff_p_int + gen_sph-&gt;ncont_fc_p_int + gen_sph-&gt;ncont_cc_p_int;

<a id="l00270" name="l00270"></a> 270 }

<a id="l00271" name="l00271"></a> 271 }();

<a id="l00272" name="l00272"></a> 272 int thread_start, thread_step, thread_stop;

<a id="l00273" name="l00273"></a> 273 Backend::ThreadIteration(thread_start, thread_step, thread_stop, n_entities);

<a id="l00274" name="l00274"></a> 274 for (int index = thread_start; index &lt; thread_stop; index += thread_step)

<a id="l00275" name="l00275"></a> 275 {

<a id="l00276" name="l00276"></a> 276 // Handling different interactionhoods structures

<a id="l00277" name="l00277"></a> 277 cuda_sph_contact* _sph_interaction = nullptr;

<a id="l00278" name="l00278"></a> 278 int n_iter, home_index, near_index, offset = 0;

<a id="l00279" name="l00279"></a> 279 if (gen_sph-&gt;use_particles_neighbors_list)

<a id="l00280" name="l00280"></a> 280 {

<a id="l00281" name="l00281"></a> 281 _sph_interaction = (it_type == itFreeCoupledOnly)

<a id="l00282" name="l00282"></a> 282 ? sph_device-&gt;cnt_sph_p + sph_device-&gt;ind_sph_p[index + gen_sph-&gt;ncont_ff_p_int]

<a id="l00283" name="l00283"></a> 283 : sph_device-&gt;cnt_sph_p + sph_device-&gt;ind_sph_p[index];

<a id="l00284" name="l00284"></a> 284 

<a id="l00285" name="l00285"></a> 285 home_index = _sph_interaction-&gt;nh;

<a id="l00286" name="l00286"></a> 286 near_index = _sph_interaction-&gt;nn;

<a id="l00287" name="l00287"></a> 287 n_iter = 1;

<a id="l00288" name="l00288"></a> 288 }

<a id="l00289" name="l00289"></a> 289 else

<a id="l00290" name="l00290"></a> 290 {

<a id="l00291" name="l00291"></a> 291 int nn_ff, nn_fc, nn_cc;

<a id="l00292" name="l00292"></a> 292 if (periodic_domain.enabled())

<a id="l00293" name="l00293"></a> 293 {

<a id="l00294" name="l00294"></a> 294 sph::get_periodic_neighbors(home_index, device_data-&gt;ind_per[index], device_data-&gt;gen, sph_device, current_time, 0.0f,

<a id="l00295" name="l00295"></a> 295 nn_ff, nn_fc, nn_cc, nn_list, nn_periodic);

<a id="l00296" name="l00296"></a> 296 if(gen_sph-&gt;nprt_active_plus_halo &gt; gen_sph-&gt;nprt_active)

<a id="l00297" name="l00297"></a> 297 sph::get_periodic_multi_GPU_neighbors(home_index, device_data-&gt;ind_per[index], device_data-&gt;gen, sph_device, current_time, 0.0f, 

<a id="l00298" name="l00298"></a> 298 nn_ff, nn_fc, nn_cc, nn_list, nn_periodic);

<a id="l00299" name="l00299"></a> 299 sph::partition_array(home_index, sph_device, nn_ff, nn_fc, nn_cc, nn_list, nn_periodic);

<a id="l00300" name="l00300"></a> 300 }

<a id="l00301" name="l00301"></a> 301 else

<a id="l00302" name="l00302"></a> 302 {

<a id="l00303" name="l00303"></a> 303 sph::get_non_periodic_neighbors(home_index, index, device_data-&gt;gen, sph_device, current_time, 0.0f, nn_ff, nn_fc, nn_cc, nn_list);

<a id="l00304" name="l00304"></a> 304 if(gen_sph-&gt;nprt_active_plus_halo &gt; gen_sph-&gt;nprt_active)

<a id="l00305" name="l00305"></a> 305 sph::get_non_periodic_multi_GPU_neighbors(home_index, device_data-&gt;gen, sph_device, current_time, 0.0f, nn_ff, nn_fc, nn_cc, nn_list);

<a id="l00306" name="l00306"></a> 306 sph::partition_array(home_index, sph_device, nn_ff, nn_fc, nn_cc, nn_list, nullptr);

<a id="l00307" name="l00307"></a> 307 }

<a id="l00308" name="l00308"></a> 308 if (it_type == itFreeAndFreeCoupled)

<a id="l00309" name="l00309"></a> 309 {

<a id="l00310" name="l00310"></a> 310 n_iter = nn_ff + nn_fc;

<a id="l00311" name="l00311"></a> 311 }

<a id="l00312" name="l00312"></a> 312 else if (it_type == itFreeCoupledOnly)

<a id="l00313" name="l00313"></a> 313 {

<a id="l00314" name="l00314"></a> 314 n_iter = nn_fc;

<a id="l00315" name="l00315"></a> 315 offset = nn_ff;

<a id="l00316" name="l00316"></a> 316 }

<a id="l00317" name="l00317"></a> 317 else // itAll

<a id="l00318" name="l00318"></a> 318 {

<a id="l00319" name="l00319"></a> 319 n_iter = nn_ff + nn_fc + nn_cc;

<a id="l00320" name="l00320"></a> 320 }

<a id="l00321" name="l00321"></a> 321 }

<a id="l00322" name="l00322"></a> 322 

<a id="l00323" name="l00323"></a> 323 // Home Interaction Interface

<a id="l00324" name="l00324"></a> 324 sph_interaction.reset_home_element(home_index);

<a id="l00325" name="l00325"></a> 325 

<a id="l00326" name="l00326"></a> 326 // Loop over number of active interactions

<a id="l00327" name="l00327"></a> 327 for (int i = 0; i &lt; n_iter; ++i)

<a id="l00328" name="l00328"></a> 328 {

<a id="l00329" name="l00329"></a> 329 // Periodic Stuff

<a id="l00330" name="l00330"></a> 330 PeriodicPosition cnt_per {};

<a id="l00331" name="l00331"></a> 331 if (!gen_sph-&gt;use_particles_neighbors_list)

<a id="l00332" name="l00332"></a> 332 {

<a id="l00333" name="l00333"></a> 333 near_index = nn_list[i + offset];

<a id="l00334" name="l00334"></a> 334 if (periodic_domain.enabled())

<a id="l00335" name="l00335"></a> 335 cnt_per = nn_periodic[i + offset];

<a id="l00336" name="l00336"></a> 336 }

<a id="l00337" name="l00337"></a> 337 else if (periodic_domain.enabled())

<a id="l00338" name="l00338"></a> 338 cnt_per = (it_type == itFreeCoupledOnly)

<a id="l00339" name="l00339"></a> 339 ? sph_device

<a id="l00340" name="l00340"></a> 340 -&gt;per_sph_cnt_p[sph_device-&gt;ind_sph_p[index + gen_sph-&gt;ncont_ff_p_int]]

<a id="l00341" name="l00341"></a> 341 : sph_device-&gt;per_sph_cnt_p[sph_device-&gt;ind_sph_p[index]];

<a id="l00342" name="l00342"></a> 342 

<a id="l00343" name="l00343"></a> 343 // Near Interaction Interface

<a id="l00344" name="l00344"></a> 344 sph_interaction.reset_near_element(near_index, cnt_per);

<a id="l00345" name="l00345"></a> 345 

<a id="l00346" name="l00346"></a> 346 // Trigger user kernel

<a id="l00347" name="l00347"></a> 347 Operator caller {};

<a id="l00348" name="l00348"></a> 348 caller(function_ptr, device_model_api, sph_model, sph_interaction_api, udf_data);

<a id="l00349" name="l00349"></a> 349 }

<a id="l00350" name="l00350"></a> 350 }

<a id="l00351" name="l00351"></a> 351}

<a id="l00352" name="l00352"></a> 352 

<a id="l00353" name="l00353"></a> 353template &lt;class Backend, typename Operator, typename Callable&gt;

<a id="l00354" name="l00354"></a> 354__host__ __device__ void launch_sph_interactions_triangles(

<a id="l00355" name="l00355"></a> 355 SDeviceModel* device_data,

<a id="l00356" name="l00356"></a> 356 SPHDeviceModel* sph_device,

<a id="l00357" name="l00357"></a> 357 double current_time,

<a id="l00358" name="l00358"></a> 358 Callable function_ptr,

<a id="l00359" name="l00359"></a> 359 void* udf_data)

<a id="l00360" name="l00360"></a> 360{

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

<a id="l00362" name="l00362"></a> 362 device_model_api._impl.model = device_data;

<a id="l00363" name="l00363"></a> 363 device_model_api._impl.current_time = current_time;

<a id="l00364" name="l00364"></a> 364 [IRockySPHDeviceModel](structIRockySPHDeviceModel.xhtml) sph_model(device_data, sph_device-&gt;gen_sph, sph_device);

<a id="l00365" name="l00365"></a> 365 

<a id="l00366" name="l00366"></a> 366 RockySPHTriangleInteraction sph_interaction(device_data, sph_device, current_time);

<a id="l00367" name="l00367"></a> 367 [IRockySPHTriangleInteraction](structIRockySPHTriangleInteraction.xhtml) sph_interaction_api { &amp;sph_interaction };

<a id="l00368" name="l00368"></a> 368 

<a id="l00369" name="l00369"></a> 369 const auto&amp; periodic_domain = device_data-&gt;gen-&gt;periodic_domain;

<a id="l00370" name="l00370"></a> 370 

<a id="l00371" name="l00371"></a> 371 // loop only over element/triangle pairs with projection inside the triangle

<a id="l00372" name="l00372"></a> 372 int thread_start, thread_step, thread_stop;

<a id="l00373" name="l00373"></a> 373 Backend::ThreadIteration(thread_start, thread_step, thread_stop, sph_device-&gt;gen_sph-&gt;n_triangle_force_pairs);

<a id="l00374" name="l00374"></a> 374 for (int i = thread_start; i &lt; thread_stop; i += thread_step)

<a id="l00375" name="l00375"></a> 375 {

<a id="l00376" name="l00376"></a> 376 const int index = sph_device-&gt;triangle_force_pairs[i];

<a id="l00377" name="l00377"></a> 377 cuda_sph_contact* _sph_interaction = sph_device-&gt;cnt_sph_t + index;

<a id="l00378" name="l00378"></a> 378 PeriodicPosition cnt_per {};

<a id="l00379" name="l00379"></a> 379 if (periodic_domain.enabled())

<a id="l00380" name="l00380"></a> 380 cnt_per = sph_device-&gt;per_sph_cnt_t[index];

<a id="l00381" name="l00381"></a> 381 

<a id="l00382" name="l00382"></a> 382 // Interaction Interface

<a id="l00383" name="l00383"></a> 383 sph_interaction.reset_triangle_contact(_sph_interaction, cnt_per);

<a id="l00384" name="l00384"></a> 384 

<a id="l00385" name="l00385"></a> 385 // ROCKY-17943: For periodic cases, skipping interaction w/ triangles, such that

<a id="l00386" name="l00386"></a> 386 // the projected point is positioned outside periodic domain

<a id="l00387" name="l00387"></a> 387 if (periodic_domain.enabled()

<a id="l00388" name="l00388"></a> 388 &amp;&amp; periodic_domain.is_outside_cartesian_periodic_bounds(sph_interaction.inside_point))

<a id="l00389" name="l00389"></a> 389 {

<a id="l00390" name="l00390"></a> 390 continue;

<a id="l00391" name="l00391"></a> 391 }

<a id="l00392" name="l00392"></a> 392 

<a id="l00393" name="l00393"></a> 393 // Trigger user kernel

<a id="l00394" name="l00394"></a> 394 Operator caller {};

<a id="l00395" name="l00395"></a> 395 caller(function_ptr, device_model_api, sph_model, sph_interaction_api, udf_data);

<a id="l00396" name="l00396"></a> 396 }

<a id="l00397" name="l00397"></a> 397}

<a id="l00398" name="l00398"></a> 398 

<a id="l00399" name="l00399"></a> 399} // namespace device

[IRockyDeviceModel](structIRockyDeviceModel.xhtml)

**Definition** rocky_model_api.hpp:19



[IRockyDeviceModel::get\_timestep](structIRockyDeviceModel.xhtml#ab03adbadfc1d2a7a8f3a1f849be73d1a)

ROCKY_FUNCTIONS double get_timestep() const

**Definition** rocky_model_api.hpp:85



[IRockyGeometryTriangle::get\_scalars](structIRockyGeometryTriangle.xhtml#a525086257224ba19ec3bbfec2522bdba)

ROCKY_FUNCTIONS IRockyTriangleScalars get_scalars()



[IRockySPHDeviceModel](structIRockySPHDeviceModel.xhtml)

**Definition** rocky_sph_model_api.hpp:146



[IRockySPHElementInteraction](structIRockySPHElementInteraction.xhtml)

**Definition** rocky_sph_interaction_api.hpp:18



[IRockySPHElementInteraction::get\_near\_element](structIRockySPHElementInteraction.xhtml#a0e9d2687e1c8a66357730003f850fe26)

ROCKY_FUNCTIONS IRockySPHElement get_near_element() const

**Definition** rocky_sph_interaction_api.hpp:112



[IRockySPHElementInteraction::add\_acceleration](structIRockySPHElementInteraction.xhtml#a713381036ee74bb8476435037c14124b)

ROCKY_FUNCTIONS void add_acceleration(float3 acceleration)

**Definition** rocky_sph_interaction_api.hpp:122



[IRockySPHElementInteraction::get\_home\_element](structIRockySPHElementInteraction.xhtml#a7eb0236ba9b25b58b0ca258b1e79976b)

ROCKY_FUNCTIONS IRockySPHElement get_home_element() const

**Definition** rocky_sph_interaction_api.hpp:107



[IRockySPHElementScalars::add\_heat\_transfer](structIRockySPHElementScalars.xhtml#ace6ee0291cca05f39eb04ba4c760afe0)

ROCKY_FUNCTIONS void add_heat_transfer(const float value)

**Definition** rocky_sph_element_scalars.hpp:243



[IRockySPHElement](structIRockySPHElement.xhtml)

**Definition** rocky_sph_element_api.hpp:106



[IRockySPHElement::get\_scalars](structIRockySPHElement.xhtml#a326bd60795936f9c4b77a9e000b1fa99)

ROCKY_FUNCTIONS IRockySPHElementScalars get_scalars()



[IRockySPHElement::add\_acceleration](structIRockySPHElement.xhtml#a8f8e5c79577f1c799383d506d2415c2b)

ROCKY_FUNCTIONS void add_acceleration(const float3 &amp;acceleration) const

**Definition** rocky_sph_element_api.hpp:146



[IRockySPHTriangleInteraction](structIRockySPHTriangleInteraction.xhtml)

**Definition** rocky_sph_interaction_api.hpp:137



[IRockySPHTriangleInteraction::add\_acceleration](structIRockySPHTriangleInteraction.xhtml#a1578dfd6da083096eb169fe24913dd5b)

ROCKY_FUNCTIONS void add_acceleration(float3 acceleration)

**Definition** rocky_sph_interaction_api.hpp:251



[IRockySPHTriangleInteraction::get\_near\_triangle](structIRockySPHTriangleInteraction.xhtml#a469b4ac53234c9e707fee62725b0470a)

ROCKY_FUNCTIONS IRockyGeometryTriangle get_near_triangle() const

**Definition** rocky_sph_interaction_api.hpp:257



[IRockySPHTriangleInteraction::get\_home\_element](structIRockySPHTriangleInteraction.xhtml#ae83d8750e222b2a46e6f901fd20b80d5)

ROCKY_FUNCTIONS IRockySPHElement get_home_element() const

**Definition** rocky_sph_interaction_api.hpp:241



[IRockyTriangleScalars::add\_scalar](structIRockyTriangleScalars.xhtml#ab9721444212bd5fe4e3af03fe3a70294)

ROCKY_FUNCTIONS void add_scalar(int scalar_index, double value) const