72 lines
1.9 KiB
Plaintext
72 lines
1.9 KiB
Plaintext
// CelerisLab – operators/helpers.cuh
|
||
// Shared device helpers used by all step kernels.
|
||
// ============================================================================
|
||
|
||
#ifndef CELERIS_OPERATORS_HELPERS_CUH
|
||
#define CELERIS_OPERATORS_HELPERS_CUH
|
||
|
||
__device__ __forceinline__ void bounce_back_swap(float* f) {
|
||
#pragma unroll
|
||
for (int i = 1; i < NQ; i += 2) {
|
||
float t = f[i]; f[i] = f[i+1]; f[i+1] = t;
|
||
}
|
||
}
|
||
|
||
__device__ __forceinline__ void collide_dispatch(
|
||
float* f, float rho_n,
|
||
#if DIM == 2
|
||
float ux, float uy
|
||
#elif DIM == 3
|
||
float ux, float uy, float uz
|
||
#endif
|
||
) {
|
||
float feq[NQ], Fin[NQ];
|
||
float omega_col = d_params.omega;
|
||
|
||
#if DIM == 2
|
||
const float fx = d_params.fx;
|
||
const float fy = d_params.fy;
|
||
if (fx != 0.0f || fy != 0.0f) {
|
||
apply_guo_velocity_correction(ux, uy, fx, fy, rho_n);
|
||
}
|
||
compute_feq(rho_n, ux, uy, feq);
|
||
if (fx != 0.0f || fy != 0.0f) {
|
||
compute_guo_forcing(ux, uy, fx, fy, Fin);
|
||
} else {
|
||
zero_forcing(Fin);
|
||
}
|
||
#elif DIM == 3
|
||
const float fx = d_params.fx;
|
||
const float fy = d_params.fy;
|
||
const float fz = d_params.fz;
|
||
if (fx != 0.0f || fy != 0.0f || fz != 0.0f) {
|
||
apply_guo_velocity_correction(ux, uy, uz, fx, fy, fz, rho_n);
|
||
}
|
||
compute_feq(rho_n, ux, uy, uz, feq);
|
||
if (fx != 0.0f || fy != 0.0f || fz != 0.0f) {
|
||
compute_guo_forcing(ux, uy, uz, fx, fy, fz, Fin);
|
||
} else {
|
||
zero_forcing(Fin);
|
||
}
|
||
#endif
|
||
|
||
#if USE_LES
|
||
omega_col = compute_omega_smag(f, feq, rho_n, omega_col);
|
||
#endif
|
||
omega_col = fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, omega_col));
|
||
|
||
#if COLLISION_MODEL == 0
|
||
collide_srt(f, feq, Fin, omega_col);
|
||
#elif COLLISION_MODEL == 1
|
||
collide_trt(f, feq, Fin, omega_col);
|
||
#elif COLLISION_MODEL == 2
|
||
#if DIM == 2
|
||
collide_mrt(f, rho_n, ux, uy, Fin, omega_col);
|
||
#elif DIM == 3
|
||
collide_mrt(f, rho_n, ux, uy, uz, Fin, omega_col);
|
||
#endif
|
||
#endif
|
||
}
|
||
|
||
#endif // CELERIS_OPERATORS_HELPERS_CUH
|