ClaudeandClaude Opus 5.5
2b3c2a7529
kernel_v2.cu -> uniform.cu (whole-domain NX/NY, Simulation) and
blocks.cu -> blocked.cu (one block window, every blocked level), with
InitTubeFlow_v2 -> InitTubeFlow. All live callers, include comments,
CONFIG.md, lifecycle doc and current-tree tests move together; no alias.
refinement.cu keeps its name until the dense-route decision.
Historical-revision tooling that names the old files (LES source pins,
timing/J campaign runners, TG run_case, supported_settings citations at
888c8038, development notes) is unchanged on purpose.
No arithmetic change; the PTX symbol InitTubeFlow_v2 is renamed.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
47 lines
1.5 KiB
Plaintext
47 lines
1.5 KiB
Plaintext
// CelerisLab -- step/esopull_macro.cu
|
|
// MacroscopicEsoPullKernel -- GPU-side macroscopic diagnostics for EsoPull.
|
|
//
|
|
// Decodes the EsoPull backing layout into physical distributions, then
|
|
// computes rho/ux/uy. Output is SoA: rho_out[N], ux_out[N], uy_out[N].
|
|
// Included by uniform.cu. No standalone compilation.
|
|
//
|
|
// This kernel is the sole path for macroscopic diagnostics in EsoPull mode.
|
|
// It uses the shared semantic helper layer (esopull_semantic_helpers.cuh).
|
|
// Do NOT attempt host-side physical DDF reconstruction for EsoPull --
|
|
// the backing layout scatters paired directions across neighbor cells.
|
|
//
|
|
// Verified: Kan99b K2 -- bit-identical to double-buffer get_macroscopic().
|
|
// ============================================================================
|
|
|
|
#ifndef CELERIS_STEP_ESOPULL_MACRO_CU
|
|
#define CELERIS_STEP_ESOPULL_MACRO_CU
|
|
|
|
__global__ void MacroscopicEsoPullKernel(
|
|
const fpxx* fi,
|
|
float* rho_out, float* ux_out, float* uy_out,
|
|
unsigned long t)
|
|
{
|
|
#if DIM == 2
|
|
unsigned int x, y; unsigned long k;
|
|
index_from_thread(x, y, k);
|
|
if (x >= (unsigned int)NX || y >= (unsigned int)NY) return;
|
|
|
|
unsigned long j[NQ];
|
|
compute_neighbors(x, y, j);
|
|
|
|
float f[NQ];
|
|
load_physical_node_esopull(k, fi, j, t, f);
|
|
|
|
float rho_n, ux, uy;
|
|
compute_rho_u(f, rho_n, ux, uy);
|
|
rho_out[k] = rho_n;
|
|
ux_out[k] = ux;
|
|
uy_out[k] = uy;
|
|
#elif DIM == 3
|
|
// TODO: D3Q19 semantic decode not implemented yet.
|
|
(void)fi; (void)rho_out; (void)ux_out; (void)uy_out; (void)t;
|
|
#endif
|
|
}
|
|
|
|
#endif // CELERIS_STEP_ESOPULL_MACRO_CU
|