223 lines
9.4 KiB
Plaintext
223 lines
9.4 KiB
Plaintext
// CelerisLab – boundary/curved_boundary.cuh
|
||
// Bouzidi linear interpolated bounce-back for curved surfaces.
|
||
//
|
||
// Pull-streaming contract used here:
|
||
// before the main pull step, each cut-link thread computes the reflected
|
||
// population and writes it into the adjacent solid source node. The pull load
|
||
// from the neighboring solid node then receives the corrected value in the
|
||
// same time step. This avoids colliding fluid nodes with garbage populations
|
||
// pulled from obstacle interiors.
|
||
//
|
||
// References:
|
||
// [Bou01] Bouzidi, Firdaouss & Lallemand, Phys. Fluids 13(11), 2001
|
||
// [Mei99] Mei, Luo & Shyy, J. Comput. Phys. 155(2), 1999
|
||
// [Gin08b] Ginzburg, Verhaeghe & d'Humieres, Commun. Comput. Phys. 2008
|
||
// [Wen15] Wen, Zhang & Fang, Entropy 17(12), 2015 (MEA force)
|
||
//
|
||
// Note on TRT:
|
||
// Plain linear Bouzidi interpolation is kept here because it is simple and
|
||
// useful for SRT and MRT baselines and for moving simple geometries.
|
||
// However, it is not a TRT-parametrized curved-wall family. Expect TRT to be
|
||
// more sensitive here than with CLI, MGLI, or other TRT-consistent schemes.
|
||
// ============================================================================
|
||
|
||
#ifndef CELERIS_BOUNDARY_CURVED_BOUNDARY_CUH
|
||
#define CELERIS_BOUNDARY_CURVED_BOUNDARY_CUH
|
||
|
||
constexpr unsigned char CURVED_FALLBACK_BOUZIDI = 0u;
|
||
constexpr unsigned char CURVED_FALLBACK_HALFWAY = 1u;
|
||
|
||
// ---------------------------------------------------------------------------
|
||
// Bouzidi linear interpolation and moving-wall correction (pure math helpers)
|
||
// ---------------------------------------------------------------------------
|
||
|
||
__device__ __forceinline__ float compute_bouzidi_reflection(
|
||
float f_toward, float f_toward_ff, float f_opp_same,
|
||
float q, unsigned char fallback_class)
|
||
{
|
||
if (fallback_class != CURVED_FALLBACK_BOUZIDI)
|
||
return f_toward;
|
||
if (q < 0.5f)
|
||
return 2.0f * q * f_toward + (1.0f - 2.0f * q) * f_toward_ff;
|
||
return (1.0f / (2.0f * q)) * f_toward +
|
||
(1.0f - 1.0f / (2.0f * q)) * f_opp_same;
|
||
}
|
||
|
||
__device__ __forceinline__ float bouzidi_linear_moving_correction(
|
||
float q, unsigned char fallback_class, float alpha_ci_dot_uw)
|
||
{
|
||
if (fallback_class != CURVED_FALLBACK_BOUZIDI) {
|
||
// Fallback replaces the q < 1/2 branch when its donor is illegal.
|
||
// Here dir points from fluid to solid and alpha_i = 3 w_i rho.
|
||
// The opposite-direction reflected population receives
|
||
// -2 * alpha_i * (c_i . u_w) in the q < 1/2 branch.
|
||
return -2.0f * alpha_ci_dot_uw;
|
||
}
|
||
if (q < 0.5f) {
|
||
return -2.0f * alpha_ci_dot_uw;
|
||
}
|
||
return -(alpha_ci_dot_uw / q);
|
||
}
|
||
|
||
__device__ __forceinline__ float compute_moving_wall_addend(
|
||
unsigned int dir, float q, unsigned char fallback_class,
|
||
float Uw, float Vw, float rho = 1.0f)
|
||
{
|
||
#if DIM == 2
|
||
const float ci_dot_uw = (float)d_cx[dir] * Uw + (float)d_cy[dir] * Vw;
|
||
#elif DIM == 3
|
||
const float ci_dot_uw = (float)d_cx[dir] * Uw + (float)d_cy[dir] * Vw + (float)d_cz[dir] * 0.0f;
|
||
#endif
|
||
const float alpha_ci_dot_uw = 3.0f * d_w[dir] * rho * ci_dot_uw;
|
||
return bouzidi_linear_moving_correction(q, fallback_class, alpha_ci_dot_uw);
|
||
}
|
||
|
||
// ---------------------------------------------------------------------------
|
||
// D2Q9: apply Bouzidi boundary condition for one cut-link (double-buffer path)
|
||
// ---------------------------------------------------------------------------
|
||
#if DIM == 2
|
||
__device__ inline void apply_bouzidi_link(
|
||
unsigned int dir, float q, unsigned long k_f,
|
||
fpxx* __restrict__ fi,
|
||
float Uw, float Vw,
|
||
float rx, float ry,
|
||
unsigned char fallback_class,
|
||
float* __restrict__ obs, int body_id, float* link_exchange = nullptr,
|
||
float* link_mass = nullptr)
|
||
{
|
||
const unsigned int dir_opp = (unsigned int)opp_dir((int)dir);
|
||
|
||
unsigned int xf, yf;
|
||
coordinates(k_f, xf, yf);
|
||
const int xs = (int)xf + d_cx[dir];
|
||
const int ys = (int)yf + d_cy[dir];
|
||
const unsigned long k_s = linear_index((unsigned int)xs, (unsigned int)ys);
|
||
|
||
const float f_toward = load_ddf(fi, index_f(k_f, dir));
|
||
|
||
float f_reflected;
|
||
if (fallback_class == CURVED_FALLBACK_BOUZIDI) {
|
||
if (q < 0.5f) {
|
||
const int xff = (int)xf - d_cx[dir];
|
||
const int yff = (int)yf - d_cy[dir];
|
||
const unsigned long k_ff = linear_index((unsigned int)xff, (unsigned int)yff);
|
||
const float f_toward_ff = load_ddf(fi, index_f(k_ff, dir));
|
||
f_reflected = compute_bouzidi_reflection(f_toward, f_toward_ff, 0.0f, q, fallback_class);
|
||
} else {
|
||
const float f_opp_same = load_ddf(fi, index_f(k_f, dir_opp));
|
||
f_reflected = compute_bouzidi_reflection(f_toward, 0.0f, f_opp_same, q, fallback_class);
|
||
}
|
||
} else {
|
||
f_reflected = compute_bouzidi_reflection(f_toward, 0.0f, 0.0f, q, fallback_class);
|
||
}
|
||
float rho_wall = USE_DDF_SHIFTING ? 1.0f : 0.0f;
|
||
if (Uw != 0.0f || Vw != 0.0f) {
|
||
for (unsigned int i = 0; i < NQ; ++i)
|
||
rho_wall += load_ddf(fi, index_f(k_f, i));
|
||
}
|
||
f_reflected += compute_moving_wall_addend(dir, q, fallback_class, Uw, Vw, rho_wall);
|
||
|
||
// Write to the solid source node so the subsequent pull step loads the
|
||
// corrected incoming population into the adjacent fluid node.
|
||
store_ddf(fi, index_f(k_s, dir_opp), f_reflected);
|
||
|
||
// Equal opposite weights cancel even with shifted DDFs.
|
||
if (link_mass != nullptr)
|
||
*link_mass = f_reflected - f_toward;
|
||
|
||
if (obs != nullptr || link_exchange != nullptr) {
|
||
#if USE_DDF_SHIFTING
|
||
const float w_i = (float)d_w[dir];
|
||
const float fx = (float)d_cx[dir] * ((f_toward + w_i) + (f_reflected + w_i));
|
||
const float fy = (float)d_cy[dir] * ((f_toward + w_i) + (f_reflected + w_i));
|
||
#else
|
||
const float fx = (float)d_cx[dir] * (f_toward + f_reflected);
|
||
const float fy = (float)d_cy[dir] * (f_toward + f_reflected);
|
||
#endif
|
||
const float tz = rx * fy - ry * fx;
|
||
if (link_exchange != nullptr) {
|
||
link_exchange[0] = fx; link_exchange[1] = fy; link_exchange[2] = tz;
|
||
}
|
||
if (obs != nullptr) {
|
||
atomicAdd(&obs[obs_force_index(body_id, 0)], fx);
|
||
atomicAdd(&obs[obs_force_index(body_id, 1)], fy);
|
||
atomicAdd(&obs[obs_torque_index(body_id, 0)], tz);
|
||
}
|
||
}
|
||
}
|
||
#endif // DIM == 2
|
||
|
||
// ---------------------------------------------------------------------------
|
||
// D3Q19: apply Bouzidi boundary condition for one cut-link (double-buffer path)
|
||
// ---------------------------------------------------------------------------
|
||
#if DIM == 3
|
||
__device__ inline void apply_bouzidi_link(
|
||
unsigned int dir, float q, unsigned long k_f,
|
||
fpxx* __restrict__ fi,
|
||
float Uw, float Vw, float Ww,
|
||
float rx, float ry, float rz,
|
||
unsigned char fallback_class,
|
||
float* __restrict__ obs, int body_id)
|
||
{
|
||
const unsigned int dir_opp = (unsigned int)opp_dir((int)dir);
|
||
|
||
unsigned int xf, yf, zf;
|
||
coordinates(k_f, xf, yf, zf);
|
||
const int xs = (int)xf + d_cx[dir];
|
||
const int ys = (int)yf + d_cy[dir];
|
||
const int zs = (int)zf + d_cz[dir];
|
||
const unsigned long k_s = linear_index((unsigned int)xs, (unsigned int)ys, (unsigned int)zs);
|
||
|
||
const float f_toward = load_ddf(fi, index_f(k_f, dir));
|
||
|
||
float f_reflected;
|
||
if (fallback_class == CURVED_FALLBACK_BOUZIDI) {
|
||
if (q < 0.5f) {
|
||
const int xff = (int)xf - d_cx[dir];
|
||
const int yff = (int)yf - d_cy[dir];
|
||
const int zff = (int)zf - d_cz[dir];
|
||
const unsigned long k_ff = linear_index((unsigned int)xff, (unsigned int)yff, (unsigned int)zff);
|
||
const float f_toward_ff = load_ddf(fi, index_f(k_ff, dir));
|
||
f_reflected = compute_bouzidi_reflection(f_toward, f_toward_ff, 0.0f, q, fallback_class);
|
||
} else {
|
||
const float f_opp_same = load_ddf(fi, index_f(k_f, dir_opp));
|
||
f_reflected = compute_bouzidi_reflection(f_toward, 0.0f, f_opp_same, q, fallback_class);
|
||
}
|
||
} else {
|
||
f_reflected = compute_bouzidi_reflection(f_toward, 0.0f, 0.0f, q, fallback_class);
|
||
}
|
||
float rho_wall = USE_DDF_SHIFTING ? 1.0f : 0.0f;
|
||
if (Uw != 0.0f || Vw != 0.0f) {
|
||
for (unsigned int i = 0; i < NQ; ++i)
|
||
rho_wall += load_ddf(fi, index_f(k_f, i));
|
||
}
|
||
f_reflected += compute_moving_wall_addend(dir, q, fallback_class, Uw, Vw, rho_wall);
|
||
|
||
store_ddf(fi, index_f(k_s, dir_opp), f_reflected);
|
||
|
||
if (obs != nullptr) {
|
||
#if USE_DDF_SHIFTING
|
||
const float w_i = (float)d_w[dir];
|
||
const float fx = (float)d_cx[dir] * ((f_toward + w_i) + (f_reflected + w_i));
|
||
const float fy = (float)d_cy[dir] * ((f_toward + w_i) + (f_reflected + w_i));
|
||
const float fz = (float)d_cz[dir] * ((f_toward + w_i) + (f_reflected + w_i));
|
||
#else
|
||
const float fx = (float)d_cx[dir] * (f_toward + f_reflected);
|
||
const float fy = (float)d_cy[dir] * (f_toward + f_reflected);
|
||
const float fz = (float)d_cz[dir] * (f_toward + f_reflected);
|
||
#endif
|
||
atomicAdd(&obs[obs_force_index(body_id, 0)], fx);
|
||
atomicAdd(&obs[obs_force_index(body_id, 1)], fy);
|
||
atomicAdd(&obs[obs_force_index(body_id, 2)], fz);
|
||
const float tx = ry * fz - rz * fy;
|
||
const float ty = rz * fx - rx * fz;
|
||
const float tz = rx * fy - ry * fx;
|
||
atomicAdd(&obs[obs_torque_index(body_id, 0)], tx);
|
||
atomicAdd(&obs[obs_torque_index(body_id, 1)], ty);
|
||
atomicAdd(&obs[obs_torque_index(body_id, 2)], tz);
|
||
}
|
||
}
|
||
#endif // DIM == 3
|
||
|
||
#endif // CELERIS_BOUNDARY_CURVED_BOUNDARY_CUH
|