## On this page

  

 

 # api\_device\_cfd\_coupling

 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// Includes =================================================================================================

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

<a id="l00005" name="l00005"></a> 5\#include &lt;rocky20/device/device\_cfd\_drag.hpp&gt;

<a id="l00006" name="l00006"></a> 6\#include &lt;rocky20/device/device\_cfd\_coupling.hpp&gt;

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

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

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

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

<a id="l00011" name="l00011"></a> 11template&lt;typename... Op&gt;

<a id="l00012" name="l00012"></a> 12__host__ __device__

<a id="l00013" name="l00013"></a> 13void execute_operators(

<a id="l00014" name="l00014"></a> 14 const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle, const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd, CFDCouplingOutput&amp; output, void* data,

<a id="l00015" name="l00015"></a> 15 Op... ops)

<a id="l00016" name="l00016"></a> 16{

<a id="l00017" name="l00017"></a> 17 int dummy[sizeof...(Op)] = { (ops(particle, cfd, output, data), 0)... };

<a id="l00018" name="l00018"></a> 18}

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

<a id="l00020" name="l00020"></a> 20/\*

<a id="l00021" name="l00021"></a> 21A utility structure to hold all device data used to execute the coupling kernel

<a id="l00022" name="l00022"></a> 22All data is a device data, so only the final kernel would be able to access such data

<a id="l00023" name="l00023"></a> 23\*/

<a id="l00024" name="l00024"></a> 24struct CouplingDispatchData

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

<a id="l00026" name="l00026"></a> 26 bool is_cuda{ false }, collect_cfd_data{ false };

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

<a id="l00028" name="l00028"></a> 28 SDeviceModel* device_data;

<a id="l00029" name="l00029"></a> 29 unsigned int* particle_indexes;

<a id="l00030" name="l00030"></a> 30 int n_particles;

<a id="l00031" name="l00031"></a> 31 void* udf_data;

<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> 35Coupling kernel 

<a id="l00036" name="l00036"></a> 36\*/

<a id="l00037" name="l00037"></a> 37template &lt;class backend, typename... Op&gt;

<a id="l00038" name="l00038"></a> 38__host__ __device__

<a id="l00039" name="l00039"></a> 39void api_coupling_kernel(CouplingDispatchData&amp; data, Op... operators)

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

<a id="l00041" name="l00041"></a> 41 auto device_data = data.device_data;

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

<a id="l00043" name="l00043"></a> 43 [IRockyParticle](structIRockyParticle.xhtml) particle_api(-1, device_data);

<a id="l00044" name="l00044"></a> 44 

<a id="l00045" name="l00045"></a> 45 [IRockyCFDProperties](structIRockyCFDProperties.xhtml) cfd{};

<a id="l00046" name="l00046"></a> 46 cfd.model = device_data;

<a id="l00047" name="l00047"></a> 47 double3 gravity { device_data-&gt;gen-&gt;gx, device_data-&gt;gen-&gt;gy, device_data-&gt;gen-&gt;gz };

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

<a id="l00049" name="l00049"></a> 49\#ifdef \_\_CUDA\_ARCH\_\_

<a id="l00050" name="l00050"></a> 50 __shared__ CFDTimestepBin block_bins[CD_NTHREADS];

<a id="l00051" name="l00051"></a> 51 const int thread_index = threadIdx.x;

<a id="l00052" name="l00052"></a> 52 CFDTimestepBin&amp; block_bin = block_bins[thread_index];

<a id="l00053" name="l00053"></a> 53\#else

<a id="l00054" name="l00054"></a> 54 CFDTimestepBin&amp; block_bin = device_data-&gt;cfd_monitor-&gt;thread_bins[backend::ThreadIndex()];

<a id="l00055" name="l00055"></a> 55\#endif

<a id="l00056" name="l00056"></a> 56 

<a id="l00057" name="l00057"></a> 57 block_bin.reset();

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

<a id="l00059" name="l00059"></a> 59 int thread_start, thread_step, thread_stop;

<a id="l00060" name="l00060"></a> 60 backend::ThreadIteration(thread_start, thread_step, thread_stop, data.n_particles);

<a id="l00061" name="l00061"></a> 61 for (auto i = thread_start; i &lt; thread_stop; i += thread_step)

<a id="l00062" name="l00062"></a> 62 {

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

<a id="l00064" name="l00064"></a> 64 const int ip = data.particle_indexes[i];

<a id="l00065" name="l00065"></a> 65 particle_api._impl.reset(ip);

<a id="l00066" name="l00066"></a> 66 if (!particle_api._impl.particle-&gt;enabled()) {

<a id="l00067" name="l00067"></a> 67 continue;

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

<a id="l00069" name="l00069"></a> 69 

<a id="l00070" name="l00070"></a> 70 const auto ic = device_data-&gt;particle_scalars-&gt;get_cfd_index(ip);

<a id="l00071" name="l00071"></a> 71 if (ic &lt; 0) {

<a id="l00072" name="l00072"></a> 72 continue;

<a id="l00073" name="l00073"></a> 73 }

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

<a id="l00075" name="l00075"></a> 75 cfd.reset(ip, ic);

<a id="l00076" name="l00076"></a> 76 CFDCouplingOutput output{};

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

<a id="l00078" name="l00078"></a> 78 // Relative Buoyancy force, due to the relative fluid density 

<a id="l00079" name="l00079"></a> 79 output.set_relative_buoyancy_force(

<a id="l00080" name="l00080"></a> 80 -cfd.get_particle_volume() * cfd.model-&gt;gen_cfd-&gt;relative_density * gravity

<a id="l00081" name="l00081"></a> 81 );

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

<a id="l00083" name="l00083"></a> 83 // Force due to pressure gradient (accounts for buoyancy and flow-induced

<a id="l00084" name="l00084"></a> 84 // pressure variations).

<a id="l00085" name="l00085"></a> 85 output.set_pressure_gradient_force(

<a id="l00086" name="l00086"></a> 86 -cfd.get_particle_volume() * cfd.get_fluid_pressure_gradient()

<a id="l00087" name="l00087"></a> 87 );

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

<a id="l00089" name="l00089"></a> 89 execute_operators(particle_api, cfd, output, data.udf_data, operators...);

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

<a id="l00091" name="l00091"></a> 91 // add forces and moments to the particle

<a id="l00092" name="l00092"></a> 92 auto&amp; particle_forces = device_data-&gt;frc[ip];

<a id="l00093" name="l00093"></a> 93 particle_forces.fx += output.total_cfd_force.x;

<a id="l00094" name="l00094"></a> 94 particle_forces.fy += output.total_cfd_force.y;

<a id="l00095" name="l00095"></a> 95 particle_forces.fz += output.total_cfd_force.z;

<a id="l00096" name="l00096"></a> 96 

<a id="l00097" name="l00097"></a> 97 if (output.has_moment()) {

<a id="l00098" name="l00098"></a> 98 particle_forces.mx += output.total_cfd_moment.x;

<a id="l00099" name="l00099"></a> 99 particle_forces.my += output.total_cfd_moment.y;

<a id="l00100" name="l00100"></a> 100 particle_forces.mz += output.total_cfd_moment.z;

<a id="l00101" name="l00101"></a> 101 }

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

<a id="l00103" name="l00103"></a> 103 // add critical time-step ratio to CFD Monitor

<a id="l00104" name="l00104"></a> 104 block_bin.check_critical_timestep(

<a id="l00105" name="l00105"></a> 105 device_data-&gt;get_general().dt,

<a id="l00106" name="l00106"></a> 106 output.drag_coefficient_fluent / particle_api.get_mass(),

<a id="l00107" name="l00107"></a> 107 device_data-&gt;cfd_monitor-&gt;levels

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

<a id="l00109" name="l00109"></a> 109 }

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

<a id="l00111" name="l00111"></a> 111\#ifdef \_\_CUDA\_ARCH\_\_

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

<a id="l00113" name="l00113"></a> 113 device::cuda_block_reduce(

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

<a id="l00115" name="l00115"></a> 115 [](CFDTimestepBin &amp;left, const CFDTimestepBin&amp; right)

<a id="l00116" name="l00116"></a> 116 {

<a id="l00117" name="l00117"></a> 117 left.join(right);

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

<a id="l00119" name="l00119"></a> 119 );

<a id="l00120" name="l00120"></a> 120\#endif

<a id="l00121" name="l00121"></a> 121}

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

<a id="l00123" name="l00123"></a> 123struct EmptyCustomCouplingPlugin {};

<a id="l00124" name="l00124"></a> 124 

<a id="l00125" name="l00125"></a> 125struct InvalidOperator

<a id="l00126" name="l00126"></a> 126{

<a id="l00127" name="l00127"></a> 127 InvalidOperator()

<a id="l00128" name="l00128"></a> 128 {

<a id="l00129" name="l00129"></a> 129 #ifndef \_\_CUDA\_ARCH\_\_

<a id="l00130" name="l00130"></a> 130 ROCKY_RUNTIME_ERROR("Invalid operation selection");

<a id="l00131" name="l00131"></a> 131 #endif

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

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

<a id="l00134" name="l00134"></a> 134 __host__ __device__ inline

<a id="l00135" name="l00135"></a> 135 void operator()(const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle, const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd, CFDCouplingOutput&amp; output, void* data)

<a id="l00136" name="l00136"></a> 136 {}

<a id="l00137" name="l00137"></a> 137};

<a id="l00138" name="l00138"></a> 138 

<a id="l00139" name="l00139"></a> 139 

<a id="l00140" name="l00140"></a> 140template &lt;class Plugin&gt;

<a id="l00141" name="l00141"></a> 141auto custom_drag_operator() { return InvalidOperator(); }

<a id="l00142" name="l00142"></a> 142 

<a id="l00143" name="l00143"></a> 143template &lt;class Plugin&gt;

<a id="l00144" name="l00144"></a> 144auto custom_lift_operator() { return InvalidOperator(); }

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

<a id="l00146" name="l00146"></a> 146template &lt;class Plugin&gt;

<a id="l00147" name="l00147"></a> 147auto custom_virtual_mass_operator(SDeviceModel *device_data) { return InvalidOperator(); }

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

<a id="l00149" name="l00149"></a> 149template &lt;class Plugin&gt;

<a id="l00150" name="l00150"></a> 150auto custom_torque_operator() { return InvalidOperator(); }

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

<a id="l00152" name="l00152"></a> 152template &lt;class Plugin&gt;

<a id="l00153" name="l00153"></a> 153auto custom_convection_operator(SDeviceModel* device_data) { return InvalidOperator(); }

<a id="l00154" name="l00154"></a> 154 

<a id="l00155" name="l00155"></a> 155/\*

<a id="l00156" name="l00156"></a> 156Coupling Operators: executed, when enabled in the following order:

<a id="l00157" name="l00157"></a> 157 - Drag

<a id="l00158" name="l00158"></a> 158 - Lift

<a id="l00159" name="l00159"></a> 159 - Virtual Mass

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

<a id="l00161" name="l00161"></a> 161 - Heat Convection

<a id="l00162" name="l00162"></a> 162 - CFD Data recorder

<a id="l00163" name="l00163"></a> 163 - Two Way Coupling

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

<a id="l00165" name="l00165"></a> 165\*/

<a id="l00166" name="l00166"></a> 166struct DragOperator

<a id="l00167" name="l00167"></a> 167{

<a id="l00168" name="l00168"></a> 168 __host__ __device__ inline

<a id="l00169" name="l00169"></a> 169 void operator()(const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle, const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd, CFDCouplingOutput&amp; output, void* data)

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

<a id="l00171" name="l00171"></a> 171 device::compute_cfd_drag(particle, cfd, output, data);

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

<a id="l00173" name="l00173"></a> 173};

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

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

<a id="l00176" name="l00176"></a> 176 

<a id="l00177" name="l00177"></a> 177template &lt;typename _DragLawOperator&gt;

<a id="l00178" name="l00178"></a> 178struct DragParametricOperator

<a id="l00179" name="l00179"></a> 179{

<a id="l00180" name="l00180"></a> 180 __host__ __device__ inline

<a id="l00181" name="l00181"></a> 181 void operator()(const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle, const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd, CFDCouplingOutput&amp; output, void* data)

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

<a id="l00183" name="l00183"></a> 183 DragLaw&lt;_DragLawOperator&gt;(particle, cfd, output, data);

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

<a id="l00185" name="l00185"></a> 185};

<a id="l00186" name="l00186"></a> 186 

<a id="l00187" name="l00187"></a> 187// Operator defined with drag laws implementation with light weight code, with less impact of performance.

<a id="l00188" name="l00188"></a> 188struct DragLightWeightOperator

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

<a id="l00190" name="l00190"></a> 190 static bool valid(const ECFDDragLaw drag_law)

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

<a id="l00192" name="l00192"></a> 192 switch (drag_law)

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

<a id="l00194" name="l00194"></a> 194 case cfd_drag_law_SchillerNaumann1933:

<a id="l00195" name="l00195"></a> 195 case cfd_drag_law_Dallavalle1948:

<a id="l00196" name="l00196"></a> 196 case cfd_drag_law_HaiderLevenspiel1989:

<a id="l00197" name="l00197"></a> 197 case cfd_drag_law_MorsiAlexander1972:

<a id="l00198" name="l00198"></a> 198 case cfd_drag_law_WenYu1966:

<a id="l00199" name="l00199"></a> 199 case cfd_drag_law_Ergun1958:

<a id="l00200" name="l00200"></a> 200 case cfd_drag_law_GidaspowBezburuahDing1992:

<a id="l00201" name="l00201"></a> 201 case cfd_drag_law_HuilinGidaspow2003:

<a id="l00202" name="l00202"></a> 202 case cfd_drag_law_DiFelice1994:

<a id="l00203" name="l00203"></a> 203 case cfd_drag_law_HillKochLadd2001:

<a id="l00204" name="l00204"></a> 204 case cfd_drag_law_HillKoch2001:

<a id="l00205" name="l00205"></a> 205 case cfd_drag_law_SyamlalOBrien1987:

<a id="l00206" name="l00206"></a> 206 return true;

<a id="l00207" name="l00207"></a> 207 default:

<a id="l00208" name="l00208"></a> 208 return false;

<a id="l00209" name="l00209"></a> 209 }

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

<a id="l00211" name="l00211"></a> 211 

<a id="l00212" name="l00212"></a> 212 __host__ __device__ inline

<a id="l00213" name="l00213"></a> 213 void operator()(const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle, const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd, CFDCouplingOutput&amp; output, void* data)

<a id="l00214" name="l00214"></a> 214 {

<a id="l00215" name="l00215"></a> 215 switch (cfd.cfd_particle_group-&gt;drag_law)

<a id="l00216" name="l00216"></a> 216 {

<a id="l00217" name="l00217"></a> 217 case cfd_drag_law_SchillerNaumann1933: DragLaw&lt;SchillerNaumann1933Operator&gt;(particle, cfd, output, data); break;

<a id="l00218" name="l00218"></a> 218 case cfd_drag_law_Dallavalle1948: DragLaw&lt;Dallavalle1948Operator&gt;(particle, cfd, output, data); break;

<a id="l00219" name="l00219"></a> 219 case cfd_drag_law_HaiderLevenspiel1989: DragLaw&lt;HaiderLevenspiel1989Operator&gt;(particle, cfd, output, data); break;

<a id="l00220" name="l00220"></a> 220 case cfd_drag_law_MorsiAlexander1972: DragLaw&lt;MorsiAlexander1972Operator&gt;(particle, cfd, output, data); break;

<a id="l00221" name="l00221"></a> 221 case cfd_drag_law_WenYu1966: DragLaw&lt;WenYu1966Operator&gt;(particle, cfd, output, data); break;

<a id="l00222" name="l00222"></a> 222 case cfd_drag_law_Ergun1958: DragLaw&lt;Ergun1958Operator&gt;(particle, cfd, output, data); break;

<a id="l00223" name="l00223"></a> 223 case cfd_drag_law_GidaspowBezburuahDing1992: DragLaw&lt;GidaspowBezburuahDing1992Operator&gt;(particle, cfd, output, data); break;

<a id="l00224" name="l00224"></a> 224 case cfd_drag_law_HuilinGidaspow2003: DragLaw&lt;HuilinGidaspow2003Operator&gt;(particle, cfd, output, data); break;

<a id="l00225" name="l00225"></a> 225 case cfd_drag_law_DiFelice1994: DragLaw&lt;DiFelice1994Operator&gt;(particle, cfd, output, data); break;

<a id="l00226" name="l00226"></a> 226 case cfd_drag_law_HillKochLadd2001: DragLaw&lt;HillKochLadd2001Operator&gt;(particle, cfd, output, data); break;

<a id="l00227" name="l00227"></a> 227 case cfd_drag_law_HillKoch2001: DragLaw&lt;HillKoch2001Operator&gt;(particle, cfd, output, data); break;

<a id="l00228" name="l00228"></a> 228 case cfd_drag_law_SyamlalOBrien1987: DragLaw&lt;SyamlalOBrien1987Operator&gt;(particle, cfd, output, data); break;

<a id="l00229" name="l00229"></a> 229 }

<a id="l00230" name="l00230"></a> 230 }

<a id="l00231" name="l00231"></a> 231};

<a id="l00232" name="l00232"></a> 232 

<a id="l00233" name="l00233"></a> 233// Operator defined with drag laws implementation with light weight code, with less impact of performance.

<a id="l00234" name="l00234"></a> 234struct DragHeavyWeightOperator

<a id="l00235" name="l00235"></a> 235{

<a id="l00236" name="l00236"></a> 236 static bool valid(const ECFDDragLaw drag_law)

<a id="l00237" name="l00237"></a> 237 {

<a id="l00238" name="l00238"></a> 238 switch (drag_law)

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

<a id="l00240" name="l00240"></a> 240 case cfd_drag_law_Ganser1993:

<a id="l00241" name="l00241"></a> 241 case cfd_drag_law_MarheinekeWegener2011:

<a id="l00242" name="l00242"></a> 242 return true;

<a id="l00243" name="l00243"></a> 243 default:

<a id="l00244" name="l00244"></a> 244 return false;

<a id="l00245" name="l00245"></a> 245 }

<a id="l00246" name="l00246"></a> 246 }

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

<a id="l00248" name="l00248"></a> 248 __host__ __device__ inline

<a id="l00249" name="l00249"></a> 249 void operator()(const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle, const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd, CFDCouplingOutput&amp; output, void* data)

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

<a id="l00251" name="l00251"></a> 251 switch (cfd.cfd_particle_group-&gt;drag_law)

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

<a id="l00253" name="l00253"></a> 253 case cfd_drag_law_Ganser1993: Ganser1993Operator::DragLaw(particle, cfd, output); break;

<a id="l00254" name="l00254"></a> 254 case cfd_drag_law_MarheinekeWegener2011: MarheinekeWegener2011Operator::DragLaw(particle, cfd, output); break;

<a id="l00255" name="l00255"></a> 255 }

<a id="l00256" name="l00256"></a> 256 }

<a id="l00257" name="l00257"></a> 257};

<a id="l00258" name="l00258"></a> 258 

<a id="l00259" name="l00259"></a> 259struct LiftOperator

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

<a id="l00261" name="l00261"></a> 261 __host__ __device__ inline

<a id="l00262" name="l00262"></a> 262 void operator()(const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle, const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd, CFDCouplingOutput&amp; output, void* data)

<a id="l00263" name="l00263"></a> 263 {

<a id="l00264" name="l00264"></a> 264 device::compute_cfd_lift(particle, cfd, output, data);

<a id="l00265" name="l00265"></a> 265 }

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

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

<a id="l00268" name="l00268"></a> 268template &lt;typename _LiftCoeff&gt;

<a id="l00269" name="l00269"></a> 269struct LiftParametricOperator

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

<a id="l00271" name="l00271"></a> 271 __host__ __device__ inline

<a id="l00272" name="l00272"></a> 272 void operator()(const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle, const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd, CFDCouplingOutput&amp; output, void* data)

<a id="l00273" name="l00273"></a> 273 {

<a id="l00274" name="l00274"></a> 274 device::LiftLaw&lt;_LiftCoeff&gt;(particle, cfd, output, data);

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

<a id="l00276" name="l00276"></a> 276};

<a id="l00277" name="l00277"></a> 277 

<a id="l00278" name="l00278"></a> 278struct VirtualMassOperator

<a id="l00279" name="l00279"></a> 279{

<a id="l00280" name="l00280"></a> 280 SDeviceModel* device_data;

<a id="l00281" name="l00281"></a> 281 

<a id="l00282" name="l00282"></a> 282 __host__ __device__ inline

<a id="l00283" name="l00283"></a> 283 void operator()(const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle, const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd, CFDCouplingOutput&amp; output, void* data)

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

<a id="l00285" name="l00285"></a> 285 device::compute_cfd_virtual_mass(particle, cfd, output, data);

<a id="l00286" name="l00286"></a> 286 device_data-&gt;particle_scalars-&gt;get_virtual_mass(particle._impl.particle_index) = output.virtual_mass;

<a id="l00287" name="l00287"></a> 287 }

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

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

<a id="l00290" name="l00290"></a> 290template &lt;typename _Coeff&gt;

<a id="l00291" name="l00291"></a> 291struct VirtualMassParametricOperator

<a id="l00292" name="l00292"></a> 292{

<a id="l00293" name="l00293"></a> 293 SDeviceModel* device_data;

<a id="l00294" name="l00294"></a> 294 

<a id="l00295" name="l00295"></a> 295 __host__ __device__ inline

<a id="l00296" name="l00296"></a> 296 void operator()(const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle, const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd, CFDCouplingOutput&amp; output, void* data)

<a id="l00297" name="l00297"></a> 297 {

<a id="l00298" name="l00298"></a> 298 device::VirtualMassLaw&lt;_Coeff&gt;(particle, cfd, output, data);

<a id="l00299" name="l00299"></a> 299 device_data-&gt;particle_scalars-&gt;get_virtual_mass(particle._impl.particle_index) = output.virtual_mass;

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

<a id="l00301" name="l00301"></a> 301};

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

<a id="l00303" name="l00303"></a> 303struct TorqueOperator

<a id="l00304" name="l00304"></a> 304{

<a id="l00305" name="l00305"></a> 305 __host__ __device__ inline

<a id="l00306" name="l00306"></a> 306 void operator()(const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle, const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd, CFDCouplingOutput&amp; output, void* data)

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

<a id="l00308" name="l00308"></a> 308 device::compute_cfd_torque(particle, cfd, output, data);

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

<a id="l00310" name="l00310"></a> 310};

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

<a id="l00312" name="l00312"></a> 312template &lt;typename _Coeff&gt;

<a id="l00313" name="l00313"></a> 313struct TorqueParametricOperator

<a id="l00314" name="l00314"></a> 314{

<a id="l00315" name="l00315"></a> 315 __host__ __device__ inline

<a id="l00316" name="l00316"></a> 316 void operator()(const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle, const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd, CFDCouplingOutput&amp; output, void* data)

<a id="l00317" name="l00317"></a> 317 {

<a id="l00318" name="l00318"></a> 318 device::TorqueLaw&lt;_Coeff&gt;(particle, cfd, output, data);

<a id="l00319" name="l00319"></a> 319 }

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

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

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

<a id="l00323" name="l00323"></a> 323{

<a id="l00324" name="l00324"></a> 324 SDeviceModel* device_data;

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

<a id="l00326" name="l00326"></a> 326 __host__ __device__ inline

<a id="l00327" name="l00327"></a> 327 void operator()(const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle, const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd, CFDCouplingOutput&amp; output, void* data)

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

<a id="l00329" name="l00329"></a> 329 device::compute_cfd_convection(particle, cfd, output, data);

<a id="l00330" name="l00330"></a> 330 device_data-&gt;particle_transfer_scalars-&gt;get_heat_transfer(particle._impl.particle_index) += output.heat_transfer;

<a id="l00331" name="l00331"></a> 331 }

<a id="l00332" name="l00332"></a> 332};

<a id="l00333" name="l00333"></a> 333 

<a id="l00334" name="l00334"></a> 334template &lt;typename _Coeff&gt;

<a id="l00335" name="l00335"></a> 335struct HeatConvectionParametricOperator

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

<a id="l00337" name="l00337"></a> 337 SDeviceModel* device_data;

<a id="l00338" name="l00338"></a> 338 

<a id="l00339" name="l00339"></a> 339 __host__ __device__ inline

<a id="l00340" name="l00340"></a> 340 void operator()(const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle, const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd, CFDCouplingOutput&amp; output, void* data)

<a id="l00341" name="l00341"></a> 341 {

<a id="l00342" name="l00342"></a> 342 device::ConvectionLaw&lt;_Coeff&gt;(particle, cfd, output, data);

<a id="l00343" name="l00343"></a> 343 device_data-&gt;particle_transfer_scalars-&gt;get_heat_transfer(particle._impl.particle_index) += output.heat_transfer;

<a id="l00344" name="l00344"></a> 344 }

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

<a id="l00346" name="l00346"></a> 346 

<a id="l00347" name="l00347"></a> 347struct CFDDataRecorderOperator

<a id="l00348" name="l00348"></a> 348{

<a id="l00349" name="l00349"></a> 349 SDeviceModel* device_data;

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

<a id="l00351" name="l00351"></a> 351 __host__ __device__ inline

<a id="l00352" name="l00352"></a> 352 void operator()(const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle, const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd, CFDCouplingOutput&amp; output, void* data)

<a id="l00353" name="l00353"></a> 353 {

<a id="l00354" name="l00354"></a> 354\#ifndef CFD\_NO\_CUSTOM\_LAWS

<a id="l00355" name="l00355"></a> 355 if (device_data-&gt;cfd_data_recorder != nullptr &amp;&amp; device_data-&gt;cfd_data_recorder-&gt;enabled)

<a id="l00356" name="l00356"></a> 356\#endif

<a id="l00357" name="l00357"></a> 357 device_data-&gt;cfd_data_recorder-&gt;register_coupling_data(

<a id="l00358" name="l00358"></a> 358 output, device_data-&gt;particle_scalars, particle._impl.particle_index, device_data-&gt;gen-&gt;dt

<a id="l00359" name="l00359"></a> 359 );

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

<a id="l00361" name="l00361"></a> 361};

<a id="l00362" name="l00362"></a> 362 

<a id="l00363" name="l00363"></a> 363struct TwoWayCouplingOperator

<a id="l00364" name="l00364"></a> 364{

<a id="l00365" name="l00365"></a> 365 SDeviceModel* device_data;

<a id="l00366" name="l00366"></a> 366 

<a id="l00367" name="l00367"></a> 367 __host__ __device__ inline void operator()(

<a id="l00368" name="l00368"></a> 368 const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle,

<a id="l00369" name="l00369"></a> 369 const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd,

<a id="l00370" name="l00370"></a> 370 CFDCouplingOutput&amp; output,

<a id="l00371" name="l00371"></a> 371 void* data)

<a id="l00372" name="l00372"></a> 372 {

<a id="l00373" name="l00373"></a> 373 if (device_data-&gt;gen_cfd-&gt;cfd_coupling_mode == cfdcTwoWay)

<a id="l00374" name="l00374"></a> 374 {

<a id="l00375" name="l00375"></a> 375 auto dt = device_data-&gt;get_general().dt;

<a id="l00376" name="l00376"></a> 376 auto particle_scalars = particle.[get\_scalars](structIRockyParticle.xhtml#a305fe472ba945e793d450053dc3b092f)();

<a id="l00377" name="l00377"></a> 377 

<a id="l00378" name="l00378"></a> 378 const auto&amp; particle_velocity = particle.[get\_translational\_velocity](structIRockyParticle.xhtml#a38420bd69fec1d3a803cbc5d79ea0197)();

<a id="l00379" name="l00379"></a> 379 const auto&amp; relative_velocity = cfd.[get\_relative\_velocity](structIRockyCFDProperties.xhtml#a891e79c0311177f586f49ead9b9a182f)();

<a id="l00380" name="l00380"></a> 380 

<a id="l00381" name="l00381"></a> 381 if (device_data-&gt;gen_cfd-&gt;use_2023R2_cfd_sources)

<a id="l00382" name="l00382"></a> 382 {

<a id="l00383" name="l00383"></a> 383 particle_scalars.add_explicit_fluid_momentum(

<a id="l00384" name="l00384"></a> 384 particle_velocity * output.drag_coefficient_fluent * dt);

<a id="l00385" name="l00385"></a> 385 }

<a id="l00386" name="l00386"></a> 386 else

<a id="l00387" name="l00387"></a> 387 {

<a id="l00388" name="l00388"></a> 388 particle_scalars.add_explicit_fluid_momentum(

<a id="l00389" name="l00389"></a> 389 -output.drag_coefficient_fluent * relative_velocity * dt);

<a id="l00390" name="l00390"></a> 390 }

<a id="l00391" name="l00391"></a> 391 particle_scalars.add_implicit_fluid_momentum(-output.drag_coefficient_fluent * dt);

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

<a id="l00393" name="l00393"></a> 393 if (output.has_explicit_forces())

<a id="l00394" name="l00394"></a> 394 {

<a id="l00395" name="l00395"></a> 395 particle_scalars.add_explicit_fluid_momentum(

<a id="l00396" name="l00396"></a> 396 -output.momentum_source_explicit_fluent * dt);

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

<a id="l00398" name="l00398"></a> 398 if (output.has_heat())

<a id="l00399" name="l00399"></a> 399 {

<a id="l00400" name="l00400"></a> 400 if (device_data-&gt;gen_cfd-&gt;use_2023R2_cfd_sources)

<a id="l00401" name="l00401"></a> 401 {

<a id="l00402" name="l00402"></a> 402 particle_scalars.add_explicit_fluid_heat_transfer_rate(

<a id="l00403" name="l00403"></a> 403 output.heat_source_explicit_fluent * dt);

<a id="l00404" name="l00404"></a> 404 }

<a id="l00405" name="l00405"></a> 405 else

<a id="l00406" name="l00406"></a> 406 {

<a id="l00407" name="l00407"></a> 407 particle_scalars.add_explicit_fluid_heat_transfer_rate(

<a id="l00408" name="l00408"></a> 408 (output.heat_source_explicit_fluent

<a id="l00409" name="l00409"></a> 409 + output.heat_source_implicit_fluent * cfd.[get\_fluid\_temperature](structIRockyCFDProperties.xhtml#a9efc8d5c0074cae18fcbcada620ae162)())

<a id="l00410" name="l00410"></a> 410 * dt);

<a id="l00411" name="l00411"></a> 411 }

<a id="l00412" name="l00412"></a> 412 particle_scalars.add_implicit_fluid_heat_transfer_rate(

<a id="l00413" name="l00413"></a> 413 output.heat_source_implicit_fluent * dt);

<a id="l00414" name="l00414"></a> 414 }

<a id="l00415" name="l00415"></a> 415 }

<a id="l00416" name="l00416"></a> 416 }

<a id="l00417" name="l00417"></a> 417};

<a id="l00418" name="l00418"></a> 418 

<a id="l00419" name="l00419"></a> 419struct TurbulenceOperator

<a id="l00420" name="l00420"></a> 420{

<a id="l00421" name="l00421"></a> 421 SDeviceModel* device_data;

<a id="l00422" name="l00422"></a> 422 

<a id="l00423" name="l00423"></a> 423 __host__ __device__ inline

<a id="l00424" name="l00424"></a> 424 void operator()(const [IRockyParticle](structIRockyParticle.xhtml)&amp; particle, const [IRockyCFDProperties](structIRockyCFDProperties.xhtml)&amp; cfd, CFDCouplingOutput&amp; output, void* data)

<a id="l00425" name="l00425"></a> 425 {

<a id="l00426" name="l00426"></a> 426\#ifndef CFD\_NO\_CUSTOM\_LAWS

<a id="l00427" name="l00427"></a> 427 if (device_data-&gt;gen_cfd-&gt;use_turbulent_dispersion)

<a id="l00428" name="l00428"></a> 428\#endif

<a id="l00429" name="l00429"></a> 429 update_turbulent_dispersion(

<a id="l00430" name="l00430"></a> 430 device_data-&gt;particle_scalars,

<a id="l00431" name="l00431"></a> 431 particle._impl.particle_index,

<a id="l00432" name="l00432"></a> 432 cfd.cfd_index,

<a id="l00433" name="l00433"></a> 433 device_data-&gt;get_general().dt,

<a id="l00434" name="l00434"></a> 434 device_data-&gt;prt_vel,

<a id="l00435" name="l00435"></a> 435 device_data-&gt;fluid_scalars,

<a id="l00436" name="l00436"></a> 436 cfd.props,

<a id="l00437" name="l00437"></a> 437 particle._impl.get_group(),

<a id="l00438" name="l00438"></a> 438 output.C_D,

<a id="l00439" name="l00439"></a> 439 device_data-&gt;cfd_random

<a id="l00440" name="l00440"></a> 440 );

<a id="l00441" name="l00441"></a> 441 }

<a id="l00442" name="l00442"></a> 442};

<a id="l00443" name="l00443"></a> 443 

<a id="l00444" name="l00444"></a> 444/\*

<a id="l00445" name="l00445"></a> 445Dispatch function, a sequence of selector function calls, that will at the end collected all operators in a 

<a id="l00446" name="l00446"></a> 446variadic template list, used to instantiate and dispatch a kernel.

<a id="l00447" name="l00447"></a> 447\*/

<a id="l00448" name="l00448"></a> 448 

<a id="l00449" name="l00449"></a> 449template &lt;class backend, typename... Operators&gt;

<a id="l00450" name="l00450"></a> 450__global__

<a id="l00451" name="l00451"></a> 451void cuda_coupling_kernel(CouplingDispatchData data, Operators... operators)

<a id="l00452" name="l00452"></a> 452{

<a id="l00453" name="l00453"></a> 453 api_coupling_kernel&lt;backend&gt;(data, operators...);

<a id="l00454" name="l00454"></a> 454}

<a id="l00455" name="l00455"></a> 455 

<a id="l00456" name="l00456"></a> 456template &lt;class backend, typename... Operators&gt;

<a id="l00457" name="l00457"></a> 457void omp_coupling_kernel(CouplingDispatchData &amp;data, Operators... operators)

<a id="l00458" name="l00458"></a> 458{

<a id="l00459" name="l00459"></a> 459 // shut up nvcc compile warnings when compiling device function version using OMPBackend

<a id="l00460" name="l00460"></a> 460 #ifndef \_\_CUDA\_ARCH\_\_

<a id="l00461" name="l00461"></a> 461 

<a id="l00462" name="l00462"></a> 462 api_coupling_kernel&lt;backend&gt;(data, operators...);

<a id="l00463" name="l00463"></a> 463 

<a id="l00464" name="l00464"></a> 464 #endif

<a id="l00465" name="l00465"></a> 465}

<a id="l00466" name="l00466"></a> 466 

<a id="l00467" name="l00467"></a> 467template &lt;typename... Operators&gt; inline

<a id="l00468" name="l00468"></a> 468void dispatch_coupling_kernel(CouplingDispatchData &amp;data, Operators... operators)

<a id="l00469" name="l00469"></a> 469{

<a id="l00470" name="l00470"></a> 470 if (data.is_cuda) {

<a id="l00471" name="l00471"></a> 471\#ifdef \_\_CUDACC\_\_

<a id="l00472" name="l00472"></a> 472 cuda_coupling_kernel&lt;CudaBackend&gt;&lt;&lt;&lt;CD_NBLOCKS, CD_NTHREADS&gt;&gt;&gt;(data, operators...);

<a id="l00473" name="l00473"></a> 473\#else

<a id="l00474" name="l00474"></a> 474 ROCKY_RUNTIME_ERROR("Invalid GPU kernel execution.");

<a id="l00475" name="l00475"></a> 475\#endif

<a id="l00476" name="l00476"></a> 476 }

<a id="l00477" name="l00477"></a> 477 else

<a id="l00478" name="l00478"></a> 478 omp_coupling_kernel&lt;OMPBackend&gt;(data, operators...);

<a id="l00479" name="l00479"></a> 479}

<a id="l00480" name="l00480"></a> 480 

<a id="l00481" name="l00481"></a> 481/\*

<a id="l00482" name="l00482"></a> 482Selector functions, called in ascending order

<a id="l00483" name="l00483"></a> 483e.g. drag -&gt; lift -&gt; virtual mass -&gt; ....

<a id="l00484" name="l00484"></a> 484\*/

<a id="l00485" name="l00485"></a> 485template&lt;typename... Operators&gt;

<a id="l00486" name="l00486"></a> 486void select_turbulence_and(CouplingDispatchData&amp; data, const cuda_general_cfd&amp; h_gen_cfd, Operators... operators)

<a id="l00487" name="l00487"></a> 487{

<a id="l00488" name="l00488"></a> 488 if (h_gen_cfd.use_turbulent_dispersion)

<a id="l00489" name="l00489"></a> 489 dispatch_coupling_kernel(data, operators..., TurbulenceOperator{ data.device_data });

<a id="l00490" name="l00490"></a> 490 else

<a id="l00491" name="l00491"></a> 491 dispatch_coupling_kernel(data, operators...);

<a id="l00492" name="l00492"></a> 492}

<a id="l00493" name="l00493"></a> 493 

<a id="l00494" name="l00494"></a> 494template&lt;typename... Operators&gt;

<a id="l00495" name="l00495"></a> 495void select_coupling_mode_and(CouplingDispatchData&amp; data, const cuda_general_cfd&amp; h_gen_cfd, Operators... operators)

<a id="l00496" name="l00496"></a> 496{

<a id="l00497" name="l00497"></a> 497 if (h_gen_cfd.cfd_coupling_mode == cfdcTwoWay)

<a id="l00498" name="l00498"></a> 498 select_turbulence_and(

<a id="l00499" name="l00499"></a> 499 data, h_gen_cfd, operators..., TwoWayCouplingOperator{ data.device_data });

<a id="l00500" name="l00500"></a> 500 else

<a id="l00501" name="l00501"></a> 501 select_turbulence_and(data, h_gen_cfd, operators...);

<a id="l00502" name="l00502"></a> 502}

<a id="l00503" name="l00503"></a> 503 

<a id="l00504" name="l00504"></a> 504template&lt;typename... Operators&gt;

<a id="l00505" name="l00505"></a> 505void select_cfd_data_recorder_and(CouplingDispatchData&amp; data, const cuda_general_cfd&amp; h_gen_cfd, Operators... operators)

<a id="l00506" name="l00506"></a> 506{

<a id="l00507" name="l00507"></a> 507\#ifdef CFD\_NO\_CUSTOM\_LAWS

<a id="l00508" name="l00508"></a> 508 if (data.collect_cfd_data)

<a id="l00509" name="l00509"></a> 509 select_coupling_mode_and(data, h_gen_cfd, operators..., CFDDataRecorderOperator{ data.device_data });

<a id="l00510" name="l00510"></a> 510 else

<a id="l00511" name="l00511"></a> 511 select_coupling_mode_and(data, h_gen_cfd, operators...);

<a id="l00512" name="l00512"></a> 512\#else

<a id="l00513" name="l00513"></a> 513 // When compiling the kernel with custom laws, we must limit the number of times we instantiate the coupling kernel.

<a id="l00514" name="l00514"></a> 514 // To avoid code bloat, we create here just a single one with all remaining functions enabled

<a id="l00515" name="l00515"></a> 515 dispatch_coupling_kernel(

<a id="l00516" name="l00516"></a> 516 data,

<a id="l00517" name="l00517"></a> 517 operators...,

<a id="l00518" name="l00518"></a> 518 CFDDataRecorderOperator{ data.device_data },

<a id="l00519" name="l00519"></a> 519 TwoWayCouplingOperator{ data.device_data },

<a id="l00520" name="l00520"></a> 520 TurbulenceOperator{ data.device_data }

<a id="l00521" name="l00521"></a> 521 );

<a id="l00522" name="l00522"></a> 522\#endif

<a id="l00523" name="l00523"></a> 523}

<a id="l00524" name="l00524"></a> 524 

<a id="l00525" name="l00525"></a> 525template&lt;class Plugin, typename... Operators&gt;

<a id="l00526" name="l00526"></a> 526void select_convection_and(CouplingDispatchData&amp; data, const cuda_general_cfd&amp; h_gen_cfd, const CFDParticleGroup&amp; group, Operators... operators)

<a id="l00527" name="l00527"></a> 527{

<a id="l00528" name="l00528"></a> 528 switch (group.convection_law)

<a id="l00529" name="l00529"></a> 529 {

<a id="l00530" name="l00530"></a> 530 case cfd_convection_law_none:

<a id="l00531" name="l00531"></a> 531 select_cfd_data_recorder_and(data, h_gen_cfd, operators...);

<a id="l00532" name="l00532"></a> 532 break;

<a id="l00533" name="l00533"></a> 533\#ifndef CFD\_NO\_CUSTOM\_LAWS

<a id="l00534" name="l00534"></a> 534 case cfd_convection_law_custom:

<a id="l00535" name="l00535"></a> 535 select_cfd_data_recorder_and(data, h_gen_cfd, operators..., custom_convection_operator&lt;Plugin&gt;(data.device_data));

<a id="l00536" name="l00536"></a> 536 break;

<a id="l00537" name="l00537"></a> 537\#endif

<a id="l00538" name="l00538"></a> 538 default:

<a id="l00539" name="l00539"></a> 539 select_cfd_data_recorder_and(data, h_gen_cfd, operators..., HeatConvectionOperator{ data.device_data });

<a id="l00540" name="l00540"></a> 540 }

<a id="l00541" name="l00541"></a> 541}

<a id="l00542" name="l00542"></a> 542 

<a id="l00543" name="l00543"></a> 543template&lt;class Plugin, typename... Operators&gt;

<a id="l00544" name="l00544"></a> 544void select_torque_and(CouplingDispatchData&amp; data, const cuda_general_cfd&amp; h_gen_cfd, const CFDParticleGroup&amp; group, Operators... operators)

<a id="l00545" name="l00545"></a> 545{

<a id="l00546" name="l00546"></a> 546 switch (group.torque_law)

<a id="l00547" name="l00547"></a> 547 {

<a id="l00548" name="l00548"></a> 548 case cfd_torque_law_none:

<a id="l00549" name="l00549"></a> 549 select_convection_and&lt;Plugin&gt;(data, h_gen_cfd, group, operators...);

<a id="l00550" name="l00550"></a> 550 break;

<a id="l00551" name="l00551"></a> 551 case cfd_torque_law_DennisSinghIngham1980:

<a id="l00552" name="l00552"></a> 552 select_convection_and&lt;Plugin&gt;(data, h_gen_cfd, group, operators..., TorqueParametricOperator&lt;DennisSinghIngham1980Operator&gt;());

<a id="l00553" name="l00553"></a> 553 break;

<a id="l00554" name="l00554"></a> 554\#ifndef CFD\_NO\_CUSTOM\_LAWS

<a id="l00555" name="l00555"></a> 555 case cfd_torque_law_custom:

<a id="l00556" name="l00556"></a> 556 select_convection_and&lt;Plugin&gt;(data, h_gen_cfd, group, operators..., custom_torque_operator&lt;Plugin&gt;());

<a id="l00557" name="l00557"></a> 557 break;

<a id="l00558" name="l00558"></a> 558\#endif

<a id="l00559" name="l00559"></a> 559 }

<a id="l00560" name="l00560"></a> 560}

<a id="l00561" name="l00561"></a> 561 

<a id="l00562" name="l00562"></a> 562template&lt;class Plugin, typename... Operators&gt;

<a id="l00563" name="l00563"></a> 563void select_virtual_mass_and(CouplingDispatchData&amp; data, const cuda_general_cfd&amp; h_gen_cfd, const CFDParticleGroup&amp; group, Operators... operators)

<a id="l00564" name="l00564"></a> 564{

<a id="l00565" name="l00565"></a> 565 switch (group.virtual_mass_law)

<a id="l00566" name="l00566"></a> 566 {

<a id="l00567" name="l00567"></a> 567 case cfd_virtual_mass_law_none:

<a id="l00568" name="l00568"></a> 568 select_torque_and&lt;Plugin&gt;(data, h_gen_cfd, group, operators...);

<a id="l00569" name="l00569"></a> 569 break;

<a id="l00570" name="l00570"></a> 570\#ifndef CFD\_NO\_CUSTOM\_LAWS

<a id="l00571" name="l00571"></a> 571 case cfd_virtual_mass_law_custom:

<a id="l00572" name="l00572"></a> 572 select_torque_and&lt;Plugin&gt;(data, h_gen_cfd, group, operators..., custom_virtual_mass_operator&lt;Plugin&gt;(data.device_data));

<a id="l00573" name="l00573"></a> 573 break;

<a id="l00574" name="l00574"></a> 574\#endif

<a id="l00575" name="l00575"></a> 575 default:

<a id="l00576" name="l00576"></a> 576 select_torque_and&lt;Plugin&gt;(data, h_gen_cfd, group, operators..., VirtualMassOperator{ data.device_data });

<a id="l00577" name="l00577"></a> 577 }

<a id="l00578" name="l00578"></a> 578}

<a id="l00579" name="l00579"></a> 579 

<a id="l00580" name="l00580"></a> 580template&lt;class Plugin, typename... Operators&gt;

<a id="l00581" name="l00581"></a> 581void select_lift_and(CouplingDispatchData&amp; data, const cuda_general_cfd&amp; h_gen_cfd, const CFDParticleGroup&amp; h_group, Operators... operators)

<a id="l00582" name="l00582"></a> 582{

<a id="l00583" name="l00583"></a> 583 switch (h_group.lift_law)

<a id="l00584" name="l00584"></a> 584 {

<a id="l00585" name="l00585"></a> 585 case cfd_lift_law_none:

<a id="l00586" name="l00586"></a> 586 select_virtual_mass_and&lt;Plugin&gt;(data, h_gen_cfd, h_group, operators...);

<a id="l00587" name="l00587"></a> 587 break;

<a id="l00588" name="l00588"></a> 588 //case cfd\_lift\_law\_Mei1992:

<a id="l00589" name="l00589"></a> 589 // select\_virtual\_mass\_and&lt;Plugin&gt;(data, h\_gen\_cfd, h\_group, operators..., LiftOperator&lt;Mei1992Operator&gt;());

<a id="l00590" name="l00590"></a> 590 // break;

<a id="l00591" name="l00591"></a> 591 //case cfd\_lift\_law\_Saffman1968:

<a id="l00592" name="l00592"></a> 592 // select\_virtual\_mass\_and&lt;Plugin&gt;(data, h\_gen\_cfd, h\_group, operators..., LiftOperator&lt;Saffman1968Operator&gt;());

<a id="l00593" name="l00593"></a> 593 // break;

<a id="l00594" name="l00594"></a> 594\#ifndef CFD\_NO\_CUSTOM\_LAWS

<a id="l00595" name="l00595"></a> 595 case cfd_lift_law_custom:

<a id="l00596" name="l00596"></a> 596 select_virtual_mass_and&lt;Plugin&gt;(data, h_gen_cfd, h_group, operators..., custom_lift_operator&lt;Plugin&gt;());

<a id="l00597" name="l00597"></a> 597 break;

<a id="l00598" name="l00598"></a> 598\#endif

<a id="l00599" name="l00599"></a> 599 default:

<a id="l00600" name="l00600"></a> 600 select_virtual_mass_and&lt;Plugin&gt;(data, h_gen_cfd, h_group, operators..., LiftOperator());

<a id="l00601" name="l00601"></a> 601 }

<a id="l00602" name="l00602"></a> 602}

<a id="l00603" name="l00603"></a> 603 

<a id="l00604" name="l00604"></a> 604template &lt;class Plugin&gt;

<a id="l00605" name="l00605"></a> 605void select_drag_and(CouplingDispatchData&amp; data, const cuda_general_cfd&amp; h_gen_cfd, const CFDParticleGroup&amp; h_group)

<a id="l00606" name="l00606"></a> 606{

<a id="l00607" name="l00607"></a> 607\#ifndef CFD\_NO\_CUSTOM\_LAWS

<a id="l00608" name="l00608"></a> 608 switch (h_group.drag_law)

<a id="l00609" name="l00609"></a> 609 {

<a id="l00610" name="l00610"></a> 610 case cfd_drag_law_custom:

<a id="l00611" name="l00611"></a> 611 select_lift_and&lt;Plugin&gt;(data, h_gen_cfd, h_group, custom_drag_operator&lt;Plugin&gt;());

<a id="l00612" name="l00612"></a> 612 break;

<a id="l00613" name="l00613"></a> 613 default: {

<a id="l00614" name="l00614"></a> 614\#else

<a id="l00615" name="l00615"></a> 615 {{

<a id="l00616" name="l00616"></a> 616\#endif

<a id="l00617" name="l00617"></a> 617 if (DragLightWeightOperator::valid(h_group.drag_law))

<a id="l00618" name="l00618"></a> 618 select_lift_and&lt;Plugin&gt;(data, h_gen_cfd, h_group, DragLightWeightOperator());

<a id="l00619" name="l00619"></a> 619 else if (DragHeavyWeightOperator::valid(h_group.drag_law))

<a id="l00620" name="l00620"></a> 620 select_lift_and&lt;Plugin&gt;(data, h_gen_cfd, h_group, DragHeavyWeightOperator());

<a id="l00621" name="l00621"></a> 621 else

<a id="l00622" name="l00622"></a> 622 select_lift_and&lt;Plugin&gt;(data, h_gen_cfd, h_group, DragOperator());

<a id="l00623" name="l00623"></a> 623 }

<a id="l00624" name="l00624"></a> 624 }

<a id="l00625" name="l00625"></a> 625}

<a id="l00626" name="l00626"></a> 626 

<a id="l00627" name="l00627"></a> 627template &lt;class Plugin&gt;

<a id="l00628" name="l00628"></a> 628void select_and_dispatch_coupling_kernel(

<a id="l00629" name="l00629"></a> 629 bool is_cuda, bool collect_cfd_data,

<a id="l00630" name="l00630"></a> 630 SDeviceModel* device_data, unsigned int* particle_indexes, int n_particles, void* udf_data,

<a id="l00631" name="l00631"></a> 631 const cuda_general_cfd&amp; h_gen_cfd, const CFDParticleGroup&amp; h_group

<a id="l00632" name="l00632"></a> 632)

<a id="l00633" name="l00633"></a> 633{

<a id="l00634" name="l00634"></a> 634 CouplingDispatchData data{ is_cuda, collect_cfd_data, device_data, particle_indexes, n_particles, udf_data };

<a id="l00635" name="l00635"></a> 635 select_drag_and&lt;Plugin&gt;(data, h_gen_cfd, h_group);

<a id="l00636" name="l00636"></a> 636}

<a id="l00637" name="l00637"></a> 637 

<a id="l00638" name="l00638"></a> 638template &lt;class backend, typename _Op&gt;

<a id="l00639" name="l00639"></a> 639__host__ __device__ void api_compute_cfd_data_coupling(

<a id="l00640" name="l00640"></a> 640 SDeviceModel* device_data, double current_time, _Op op, void* udf_data)

<a id="l00641" name="l00641"></a> 641{

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

<a id="l00643" name="l00643"></a> 643 device_model_api._impl.model = device_data;

<a id="l00644" name="l00644"></a> 644 device_model_api._impl.current_time = current_time;

<a id="l00645" name="l00645"></a> 645 

<a id="l00646" name="l00646"></a> 646 [IRockyParticle](structIRockyParticle.xhtml) particle_api(-1, device_data);

<a id="l00647" name="l00647"></a> 647 

<a id="l00648" name="l00648"></a> 648 [IRockyCFDProperties](structIRockyCFDProperties.xhtml) cfd{};

<a id="l00649" name="l00649"></a> 649 cfd.model = device_data;

<a id="l00650" name="l00650"></a> 650 

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

<a id="l00652" name="l00652"></a> 652 

<a id="l00653" name="l00653"></a> 653 int thread_start, thread_step, thread_stop;

<a id="l00654" name="l00654"></a> 654 backend::ThreadIteration(thread_start, thread_step, thread_stop, general.nprt_active);

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

<a id="l00656" name="l00656"></a> 656 {

<a id="l00657" name="l00657"></a> 657 particle_api._impl.reset(i);

<a id="l00658" name="l00658"></a> 658 int ind_cfd = device_data-&gt;particle_scalars-&gt;get_cfd_index(i);

<a id="l00659" name="l00659"></a> 659 if (!particle_api._impl.particle-&gt;enabled() || ind_cfd &lt; 0)

<a id="l00660" name="l00660"></a> 660 {

<a id="l00661" name="l00661"></a> 661 continue;

<a id="l00662" name="l00662"></a> 662 }

<a id="l00663" name="l00663"></a> 663 

<a id="l00664" name="l00664"></a> 664 cfd.reset(i, ind_cfd);

<a id="l00665" name="l00665"></a> 665 

<a id="l00666" name="l00666"></a> 666 op(device_model_api, particle_api, cfd, udf_data);

<a id="l00667" name="l00667"></a> 667 }

<a id="l00668" name="l00668"></a> 668}

<a id="l00669" name="l00669"></a> 669 

<a id="l00670" name="l00670"></a> 670} // namespace device

[IRockyCFDProperties](structIRockyCFDProperties.xhtml)

**Definition** rocky_cfd.hpp:26



[IRockyCFDProperties::get\_relative\_velocity](structIRockyCFDProperties.xhtml#a891e79c0311177f586f49ead9b9a182f)

ROCKY_FUNCTIONS double3 get_relative_velocity() const

**Definition** rocky_cfd.hpp:112



[IRockyCFDProperties::get\_fluid\_temperature](structIRockyCFDProperties.xhtml#a9efc8d5c0074cae18fcbcada620ae162)

ROCKY_FUNCTIONS double get_fluid_temperature() const

**Definition** rocky_cfd.hpp:54



[IRockyDeviceModel](structIRockyDeviceModel.xhtml)

**Definition** rocky_model_api.hpp:19



[IRockyParticle](structIRockyParticle.xhtml)

**Definition** rocky_particle_api.hpp:27



[IRockyParticle::get\_scalars](structIRockyParticle.xhtml#a305fe472ba945e793d450053dc3b092f)

ROCKY_FUNCTIONS IRockyParticleScalars get_scalars()

**Definition** rocky_particle_api.hpp:537



[IRockyParticle::get\_translational\_velocity](structIRockyParticle.xhtml#a38420bd69fec1d3a803cbc5d79ea0197)

ROCKY_FUNCTIONS double3 get_translational_velocity() const

**Definition** rocky_particle_api.hpp:417