Frank14f/CelerisLab
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Add TRT-LES diagnostics and plotting functionality

Browse Source 
- Introduced new functions for computing and plotting TRT-LES fields.
- Enhanced case diagnostics to include mass drift and inlet variance metrics.
- Updated configuration to support inlet profile selection and TRT magic parameter.
- Modified existing functions to accommodate new diagnostic calculations.
- Improved case tagging to include inlet profile and TRT parameters.
- Added checks for fluid dynamics and diagnostics in the run_case function.
This commit is contained in:
Frank14f
2026-04-06 15:57:44 +08:00
parent 5076e5d789
commit b339f302d3
16 changed files with 632 additions and 773 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/CelerisLab/lbm/driver.py
+7
View File
@@ -193,12 +193,17 @@ class FlowField:
# max_id = max(self.objects.keys())
else:
raise ValueError(f"Unsupported data type {self.DATA_TYPE}.")
# Ensure host-side DDF mirrors current device state before local edits.
cuda.memcpy_dtoh(self.ddf, self.ddf_gpu)
for x in range(int(x_c - radius) - 1, int(x_c + radius) + 1):
for y in range(int(y_c - radius) - 1, int(y_c + radius) + 1):
if (x - x_c) ** 2 + (y - y_c) ** 2 < radius**2:
k = x + y * self.FIELD_SHAPE[0]
self.flag[k] = SOLID
for i in range(self.LATTICE):
self.ddf[k + i * self.FIELD_SIZE] = self.WW[i]
delta_temp = np.zeros(11, dtype=self.DATA_TYPE)
delta_temp[0] = id_object.view(self.DATA_TYPE)
for i in range(self.LATTICE):
@@ -245,6 +250,8 @@ class FlowField:
cuda.memcpy_htod(self.delta_gpu, self.delta_curve)
cuda.memcpy_htod(self.flag_gpu, self.flag)
cuda.memcpy_htod(self.indx_gpu, self.indx)
cuda.memcpy_htod(self.ddf_gpu, self.ddf)
cuda.memcpy_htod(self.temp_gpu, self.ddf)
self._rebuild_kernel()
src/CelerisLab/lbm/kernels/boundary/bounce_back.cuh
+19
View File
@@ -43,6 +43,25 @@ __device__ inline void apply_bounce_back(float* __restrict__ f,
// For nodes at y=1 (adjacent to y=0 wall) or y=NY-2 (adjacent to y=NY-1 wall)
// ---------------------------------------------------------------------------
#if NQ == 9
__device__ inline void apply_wall_bb_d2q9_y_pull(unsigned int y,
float* __restrict__ f,
const fpxx* __restrict__ fi_in,
unsigned long k)
{
if (y == 1) {
// Directions sourced from y=0 wall in pull step: +y, +x+y, -x+y.
f[3] = load_ddf(fi_in, index_f(k, 4u));
f[5] = load_ddf(fi_in, index_f(k, 6u));
f[8] = load_ddf(fi_in, index_f(k, 7u));
}
else if (y == (unsigned int)(NY - 2)) {
// Directions sourced from y=NY-1 wall in pull step: -y, -x-y, +x-y.
f[4] = load_ddf(fi_in, index_f(k, 3u));
f[6] = load_ddf(fi_in, index_f(k, 5u));
f[7] = load_ddf(fi_in, index_f(k, 8u));
}
}
__device__ inline void apply_wall_bb_d2q9(unsigned int y,
float* __restrict__ f)
{
src/CelerisLab/lbm/kernels/boundary/inlet_outlet.cuh
+49 -12
View File
@@ -31,6 +31,14 @@
#define OUTLET_BLEND_ALPHA 0.70f
#endif
#ifndef OUTLET_SRT_NEQ_DAMP
#define OUTLET_SRT_NEQ_DAMP 0.50f
#endif
#ifndef INLET_TRT_NEQ_DAMP
#define INLET_TRT_NEQ_DAMP 0.50f
#endif
__device__ __forceinline__ float inlet_target_u(float y_coord) {
#if INLET_PROFILE == 0
return U0;
@@ -69,10 +77,22 @@ __device__ inline void apply_parabolic_inlet(float* __restrict__ f,
compute_feq(rho_neb, u_target, v_target, feq_tar);
compute_feq(rho_neb, u_neb, v_neb, feq_neb);
// Non-equilibrium extrapolation
f[1] = f_neb[1] - feq_neb[1] + feq_tar[1];
f[5] = f_neb[5] - feq_neb[5] + feq_tar[5];
f[7] = f_neb[7] - feq_neb[7] + feq_tar[7];
#if COLLISION_MODEL == 1
// TRT path: reconstruct full population set at inlet using damped donor
// non-equilibrium transport. This follows the high-Re stable family that
// replaces all boundary-node populations, reducing odd-mode contamination.
const float beta = INLET_TRT_NEQ_DAMP;
#pragma unroll
for (int i = 0; i < 9; i++) {
const float fneq = f_neb[i] - feq_neb[i];
f[i] = feq_tar[i] + beta * fneq;
}
#else
const float beta = 1.0f;
f[1] = feq_tar[1] + beta * (f_neb[1] - feq_neb[1]);
f[5] = feq_tar[5] + beta * (f_neb[5] - feq_neb[5]);
f[7] = feq_tar[7] + beta * (f_neb[7] - feq_neb[7]);
#endif
}
// ---------------------------------------------------------------------------
@@ -94,20 +114,30 @@ __device__ inline void apply_pressure_outlet(float* __restrict__ f,
f[6] = f_neb[6];
#else
// Convective non-reflecting style: extrapolate outlet density from neighbor
// and keep neighbor velocity.
// Prescribed-pressure outlet: keep neighbor velocity and impose reference
// outlet density for NEQ reconstruction.
float rho_neb, u_neb, v_neb;
compute_rho_u(f_neb, rho_neb, u_neb, v_neb);
#if OUTLET_BACKFLOW_CLAMP
u_neb = fmaxf(u_neb, 0.0f);
#endif
float rho_out = rho_neb;
float rho_out = RHO;
float feq_tar[9], feq_neb[9];
compute_feq(rho_out, u_neb, v_neb, feq_tar);
compute_feq(rho_neb, u_neb, v_neb, feq_neb);
#if OUTLET_MODE == 2
#if COLLISION_MODEL == 0
// SRT path: use full-population damped NEQ reconstruction at outlet to
// suppress checkerboard/grid noise from high-frequency donor content.
// Reference: high-Re outlet regularization rationale and NEQ decomposition.
const float beta = OUTLET_SRT_NEQ_DAMP;
#pragma unroll
for (int i = 0; i < 9; i++) {
const float fneq = f_neb[i] - feq_neb[i];
f[i] = feq_tar[i] + beta * fneq;
}
#elif OUTLET_MODE == 2
const float a = OUTLET_BLEND_ALPHA;
f[2] = a * (f_neb[2] - feq_neb[2] + feq_tar[2]) + (1.0f - a) * f_neb[2];
f[8] = a * (f_neb[8] - feq_neb[8] + feq_tar[8]) + (1.0f - a) * f_neb[8];
@@ -178,15 +208,22 @@ __device__ inline void apply_pressure_outlet_3d(float* __restrict__ f,
un = fmaxf(un, 0.0f);
#endif
// Convective non-reflecting style: extrapolate outlet density from neighbor
// and keep neighbor velocity.
float rho_out = rho_neb;
// Prescribed-pressure outlet: keep neighbor velocity and impose reference
// outlet density for NEQ reconstruction.
float rho_out = RHO;
float feq_tar[19], feq_neb[19];
compute_feq(rho_out, un, vn, wn, feq_tar);
compute_feq(rho_neb, un, vn, wn, feq_neb);
// Reconstruct cx<0 directions: i = 2, 8, 10, 14, 16
#if OUTLET_MODE == 2
#if COLLISION_MODEL == 0
const float beta = OUTLET_SRT_NEQ_DAMP;
#pragma unroll
for (int i = 0; i < 19; i++) {
const float fneq = f_neb[i] - feq_neb[i];
f[i] = feq_tar[i] + beta * fneq;
}
#elif OUTLET_MODE == 2
const float a = OUTLET_BLEND_ALPHA;
f[2] = a * (f_neb[2] - feq_neb[2] + feq_tar[2]) + (1.0f - a) * f_neb[2];
f[8] = a * (f_neb[8] - feq_neb[8] + feq_tar[8]) + (1.0f - a) * f_neb[8];
src/CelerisLab/lbm/kernels/kernel_v2.cu
+14 -6
View File
@@ -91,9 +91,8 @@ __global__ void InitTubeFlow_v2(uint8_t* flag, fpxx* fi)
index_from_thread(x, y, k);
if (x >= (unsigned int)NX || y >= (unsigned int)NY) return;
float u_init = U0 * 1.5f * (1.0f - 4.0f * ((float)y - 0.5f * (NY - 1))
* ((float)y - 0.5f * (NY - 1))
/ ((float)(NY - 2) * (float)(NY - 2)));
// Keep startup field consistent with runtime inlet configuration.
float u_init = inlet_target_u((float)y);
if (y == 0 || y == NY - 1 || x == 0 || x == NX - 1) {
flag[k] = LEGACY_SOLID;
@@ -116,9 +115,8 @@ __global__ void InitTubeFlow_v2(uint8_t* flag, fpxx* fi)
index_from_thread(x, y, z, k);
if (x >= (unsigned int)NX || y >= (unsigned int)NY || z >= (unsigned int)NZ) return;
float u_init = U0 * 1.5f * (1.0f - 4.0f * ((float)y - 0.5f * (NY - 1))
* ((float)y - 0.5f * (NY - 1))
/ ((float)(NY - 2) * (float)(NY - 2)));
// Keep startup field consistent with runtime inlet configuration.
float u_init = inlet_target_u((float)y);
if (y == 0 || y == NY - 1 || x == 0 || x == NX - 1) {
flag[k] = LEGACY_SOLID;
@@ -217,9 +215,14 @@ __global__ void OneStep(
}
}
if ((fl & LEGACY_FLUID) && (y == 1u || y == (unsigned int)(NY - 2))) {
apply_wall_bb_d2q9_y_pull(y, f, fi_in, k);
}
// Collision
if (fl & LEGACY_FLUID) {
float feq[NQ], Fin[NQ];
compute_rho_u(f, rho_n, ux, uy);
compute_feq(rho_n, ux, uy, feq);
zero_forcing(Fin);
float omega_col = d_params.omega;
@@ -298,9 +301,14 @@ __global__ void OneStep(
}
}
if ((fl & LEGACY_FLUID) && (y == 1 || y == NY - 2)) {
apply_wall_bb_d3q19_y_pull(y, f, fi_in, k);
}
// ---- Collision (fluid only) ----
if (fl & LEGACY_FLUID) {
float feq[NQ], Fin[NQ];
compute_rho_u(f, rho_n, ux, uy, uz);
compute_feq(rho_n, ux, uy, uz, feq);
zero_forcing(Fin);
float omega_col = d_params.omega;
src/CelerisLab/lbm/kernels/macros.h
+5
View File
@@ -100,4 +100,9 @@
#ifndef OMEGA_COLLISION_MAX
#define OMEGA_COLLISION_MAX 1.999f
#endif
// TRT magic parameter Lambda used to map omega+ -> omega-
#ifndef TRT_MAGIC_PARAM
#define TRT_MAGIC_PARAM 0.1875f
#endif
src/CelerisLab/lbm/kernels/operators/collision_trt.cuh
+4 -1
View File
@@ -13,8 +13,11 @@
#ifndef CELERIS_OPERATORS_COLLISION_TRT_CUH
#define CELERIS_OPERATORS_COLLISION_TRT_CUH
// Magic parameter Λ = 3/16 (optimal for porous-media / bounce-back wall location)
// Magic parameter Λ for TRT wall/collision coupling.
// Default keeps previous behavior; can be overridden in macros.h / tests.
#ifndef TRT_MAGIC_PARAM
#define TRT_MAGIC_PARAM (0.1875f)
#endif
__device__ __forceinline__ float compute_omega_minus(float omega_plus) {
return 1.0f / (TRT_MAGIC_PARAM / (1.0f / omega_plus - 0.5f) + 0.5f);
src/CelerisLab/lbm/kernels/operators/turbulence_smag.cuh
+78 -25
View File
@@ -19,8 +19,21 @@
#define LES_CS 0.16f
#endif
#ifndef LES_FP_ITERS
#define LES_FP_ITERS 3
#endif
#ifndef LES_FP_RELAX
#define LES_FP_RELAX 0.70f
#endif
#ifndef LES_NUT_MAX_RATIO
#define LES_NUT_MAX_RATIO 20.0f
#endif
__device__ __forceinline__ float clamp_omega(float w) {
return fminf(1.999f, fmaxf(0.01f, w));
return fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, w));
}
#if NQ == 9
@@ -32,6 +45,7 @@ __device__ __forceinline__ float compute_omega_smag(const float* __restrict__ f,
{
const float rho_safe = fmaxf(rho, 1.0e-12f);
const float tau0 = fmaxf(1.0f / fmaxf(omega0, 1.0e-6f), 0.500001f);
const float nu0 = (tau0 - 0.5f) * (1.0f / 3.0f);
// Πneq = Σ (ci ci)(fi-feqi)
float pixx = 0.0f, piyy = 0.0f, pixy = 0.0f;
@@ -46,18 +60,36 @@ __device__ __forceinline__ float compute_omega_smag(const float* __restrict__ f,
pixy += fneq * cx * cy;
}
const float denom = 2.0f * rho_safe * CS2 * tau0;
const float sxx = -pixx / denom;
const float syy = -piyy / denom;
const float sxy = -pixy / denom;
// Self-consistent tau iteration for LES/TRT coupling.
// S is estimated from Pi_neq with current tau_eff, then nu_t updates tau_eff.
float tau_eff = tau0;
const float tau_max = 1.0f / fmaxf(OMEGA_COLLISION_MIN, 1.0e-6f);
// |S| = sqrt(2 Sij Sij)
const float s_mag = sqrtf(2.0f * (sxx*sxx + syy*syy + 2.0f*sxy*sxy));
#pragma unroll
for (int it = 0; it < LES_FP_ITERS; ++it) {
const float denom = 2.0f * rho_safe * CS2 * fmaxf(tau_eff, 0.500001f);
const float sxx = -pixx / denom;
const float syy = -piyy / denom;
const float sxy = -pixy / denom;
const float nu0 = (tau0 - 0.5f) * (1.0f / 3.0f);
const float nut = (LES_CS * LES_CS) * s_mag;
const float nue = nu0 + nut;
return clamp_omega(1.0f / (3.0f * nue + 0.5f));
// Use deviatoric strain to reduce SGS sensitivity to bulk/pressure-wave content.
const float tr = 0.5f * (sxx + syy);
const float sxx_dev = sxx - tr;
const float syy_dev = syy - tr;
const float sxy_dev = sxy;
// |S_dev| = sqrt(2 S_dev,ij S_dev,ij)
const float s_mag = sqrtf(2.0f * (sxx_dev*sxx_dev + syy_dev*syy_dev + 2.0f*sxy_dev*sxy_dev));
float nut = (LES_CS * LES_CS) * s_mag;
const float nut_cap = fmaxf(0.0f, LES_NUT_MAX_RATIO * nu0);
nut = fminf(fmaxf(nut, 0.0f), nut_cap);
const float tau_new = fminf(tau_max, fmaxf(0.500001f, 0.5f + 3.0f * (nu0 + nut)));
tau_eff = LES_FP_RELAX * tau_new + (1.0f - LES_FP_RELAX) * tau_eff;
}
return clamp_omega(1.0f / tau_eff);
}
#elif NQ == 19
@@ -69,6 +101,7 @@ __device__ __forceinline__ float compute_omega_smag(const float* __restrict__ f,
{
const float rho_safe = fmaxf(rho, 1.0e-12f);
const float tau0 = fmaxf(1.0f / fmaxf(omega0, 1.0e-6f), 0.500001f);
const float nu0 = (tau0 - 0.5f) * (1.0f / 3.0f);
// Πneq = Σ (ci ci)(fi-feqi)
float pixx = 0.0f, piyy = 0.0f, pizz = 0.0f;
@@ -88,22 +121,42 @@ __device__ __forceinline__ float compute_omega_smag(const float* __restrict__ f,
piyz += fneq * cy * cz;
}
const float denom = 2.0f * rho_safe * CS2 * tau0;
const float sxx = -pixx / denom;
const float syy = -piyy / denom;
const float szz = -pizz / denom;
const float sxy = -pixy / denom;
const float sxz = -pixz / denom;
const float syz = -piyz / denom;
// Self-consistent tau iteration for LES/TRT coupling.
float tau_eff = tau0;
const float tau_max = 1.0f / fmaxf(OMEGA_COLLISION_MIN, 1.0e-6f);
// |S| = sqrt(2 Sij Sij)
const float s_mag = sqrtf(2.0f * (sxx*sxx + syy*syy + szz*szz
+ 2.0f*(sxy*sxy + sxz*sxz + syz*syz)));
#pragma unroll
for (int it = 0; it < LES_FP_ITERS; ++it) {
const float denom = 2.0f * rho_safe * CS2 * fmaxf(tau_eff, 0.500001f);
const float sxx = -pixx / denom;
const float syy = -piyy / denom;
const float szz = -pizz / denom;
const float sxy = -pixy / denom;
const float sxz = -pixz / denom;
const float syz = -piyz / denom;
const float nu0 = (tau0 - 0.5f) * (1.0f / 3.0f);
const float nut = (LES_CS * LES_CS) * s_mag;
const float nue = nu0 + nut;
return clamp_omega(1.0f / (3.0f * nue + 0.5f));
// Use deviatoric strain to reduce SGS sensitivity to bulk/pressure-wave content.
const float tr = (sxx + syy + szz) / 3.0f;
const float sxx_dev = sxx - tr;
const float syy_dev = syy - tr;
const float szz_dev = szz - tr;
const float sxy_dev = sxy;
const float sxz_dev = sxz;
const float syz_dev = syz;
// |S_dev| = sqrt(2 S_dev,ij S_dev,ij)
const float s_mag = sqrtf(2.0f * (sxx_dev*sxx_dev + syy_dev*syy_dev + szz_dev*szz_dev
+ 2.0f*(sxy_dev*sxy_dev + sxz_dev*sxz_dev + syz_dev*syz_dev)));
float nut = (LES_CS * LES_CS) * s_mag;
const float nut_cap = fmaxf(0.0f, LES_NUT_MAX_RATIO * nu0);
nut = fminf(fmaxf(nut, 0.0f), nut_cap);
const float tau_new = fminf(tau_max, fmaxf(0.500001f, 0.5f + 3.0f * (nu0 + nut)));
tau_eff = LES_FP_RELAX * tau_new + (1.0f - LES_FP_RELAX) * tau_eff;
}
return clamp_omega(1.0f / tau_eff);
}
#endif // NQ
src/CelerisLab/lbm/kernels/step/one_step_double.cu
+6 -4
View File
@@ -128,12 +128,12 @@ __global__ void StreamCollideDouble(
// ----- Wall bounce-back (top/bottom walls) -----
#if NQ == 9
if (y == 1 || y == NY - 2) {
apply_wall_bb_d2q9(y, f);
if ((fl & LEGACY_FLUID) && (y == 1u || y == (unsigned int)(NY - 2))) {
apply_wall_bb_d2q9_y_pull(y, f, fi_in, k);
}
#elif NQ == 19
if (y == 1 || y == NY - 2) {
apply_wall_bb_d3q19_y(y, f);
if ((fl & LEGACY_FLUID) && (y == 1u || y == (unsigned int)(NY - 2))) {
apply_wall_bb_d3q19_y_pull(y, f, fi_in, k);
}
#endif
@@ -143,6 +143,7 @@ __global__ void StreamCollideDouble(
float Fin[NQ];
#if NQ == 9
compute_rho_u(f, rho_n, ux, uy);
compute_feq(rho_n, ux, uy, feq);
zero_forcing(Fin);
float omega_col = d_params.omega;
@@ -161,6 +162,7 @@ __global__ void StreamCollideDouble(
#endif
#elif NQ == 19
compute_rho_u(f, rho_n, ux, uy, uz);
compute_feq(rho_n, ux, uy, uz, feq);
zero_forcing(Fin);
float omega_col = d_params.omega;