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>
79 lines
3.6 KiB
Plaintext
79 lines
3.6 KiB
Plaintext
// CelerisLab/lbm/kernels/step/block_transfer.cu
|
|
// Same-level halo refresh and 2:1 explosion/coalescence between blocked levels.
|
|
// Part of the blocked.cu compilation entry; relies on its includes and guards.
|
|
|
|
// Halo refresh: block cell ids b*TOTAL_CELLS+c, sources are always core cells.
|
|
extern "C" __global__ void BlockHaloCopy(
|
|
fpxx* f, const uint64_t* dst, const uint64_t* src, unsigned int count)
|
|
{
|
|
unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (row >= count) return;
|
|
const uint64_t d = dst[row], s = src[row];
|
|
fpxx* to = f + (d / TOTAL_CELLS) * NQ * TOTAL_CELLS;
|
|
const fpxx* from = f + (s / TOTAL_CELLS) * NQ * TOTAL_CELLS;
|
|
for (unsigned int q = 0; q < NQ; ++q)
|
|
to[index_f(d % TOTAL_CELLS, q)] = from[index_f(s % TOTAL_CELLS, q)];
|
|
}
|
|
|
|
// Fixed-patch explosion into blocked fine storage: target = owner block cell
|
|
// id of a patch-ghost cell; donors are parent block cell ids. Dense parent
|
|
// storage is the one-window case with coarse_cells = full domain size.
|
|
extern "C" __global__ void BlockExplode(
|
|
const fpxx* coarse, fpxx* fine, const uint64_t* target, const uint64_t* donor,
|
|
unsigned int count, unsigned long coarse_cells)
|
|
{
|
|
unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (row >= count) return;
|
|
const uint64_t t = target[row];
|
|
fpxx* to = fine + (t / TOTAL_CELLS) * NQ * TOTAL_CELLS;
|
|
const uint64_t d = donor[row];
|
|
const fpxx* from = coarse + (d / coarse_cells) * NQ * coarse_cells;
|
|
for (unsigned int q = 0; q < NQ; ++q)
|
|
to[index_f(t % TOTAL_CELLS, q)] = from[q * coarse_cells + d % coarse_cells];
|
|
}
|
|
|
|
// Experimental Schornbaum--Ruede (arXiv:1508.07982, Algorithm 2)
|
|
// and Chen et al. (doi:10.1016/j.physa.2005.09.036) population interpolation.
|
|
// Independent implementation of the mathematical stencil; no upstream source copied.
|
|
// Stencil cells are owner cores; a disabled axis repeats its center twice.
|
|
extern "C" __global__ void BlockExplodeProjected(
|
|
const fpxx* coarse, fpxx* fine, const uint64_t* target, const uint64_t* donor,
|
|
unsigned int count, unsigned long coarse_cells,
|
|
const uint64_t* neighbors, const float* offsets)
|
|
{
|
|
unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (row >= count) return;
|
|
const uint64_t t = target[row], d = donor[row];
|
|
fpxx* to = fine + (t / TOTAL_CELLS) * NQ * TOTAL_CELLS;
|
|
for (unsigned int q = 0; q < NQ; ++q) {
|
|
float v[4];
|
|
for (unsigned int a = 0; a < 4; ++a) {
|
|
const uint64_t n = neighbors[4ul * row + a];
|
|
v[a] = coarse[(n / coarse_cells) * NQ * coarse_cells
|
|
+ q * coarse_cells + n % coarse_cells];
|
|
}
|
|
float gx = .5f * (v[1] - v[0]), gy = .5f * (v[3] - v[2]);
|
|
const float ex = d_cx[q], ey = d_cy[q];
|
|
if (q != 0) {
|
|
const float along = (ex * gx + ey * gy) / (ex * ex + ey * ey);
|
|
gx -= ex * along;
|
|
gy -= ey * along;
|
|
}
|
|
to[index_f(t % TOTAL_CELLS, q)] =
|
|
coarse[(d / coarse_cells) * NQ * coarse_cells + q * coarse_cells + d % coarse_cells]
|
|
+ offsets[2ul * row] * gx + offsets[2ul * row + 1] * gy;
|
|
}
|
|
}
|
|
|
|
// Targets are receiver-local SOURCE slots for the immediately following pull.
|
|
// Keep child summation order identical to PatchCoarseStep; no collision here.
|
|
extern "C" __global__ void BlockCoalesce(
|
|
const fpxx* fine, fpxx* coarse, const uint64_t* target, const uint64_t* children,
|
|
unsigned int count)
|
|
{
|
|
unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
|
|
if (row >= count) return;
|
|
const uint64_t* s = children + 4ul * row;
|
|
coarse[target[row]] = 0.25f * (fine[s[0]] + fine[s[1]] + fine[s[2]] + fine[s[3]]);
|
|
}
|