diff --git a/crates/specialized/rtx-cfd/src/kernels/cuda/mg_vcycle.cu b/crates/specialized/rtx-cfd/src/kernels/cuda/mg_vcycle.cu index 601c872..837c9dd 100644 --- a/crates/specialized/rtx-cfd/src/kernels/cuda/mg_vcycle.cu +++ b/crates/specialized/rtx-cfd/src/kernels/cuda/mg_vcycle.cu @@ -78,9 +78,13 @@ extern "C" __global__ void mg_zero(int n_cells, const unsigned int* __restrict__ x[(size_t)blockIdx.y * n + cells[t]] = 0.0f; } -/* The coarsest level: one thread per march, `sweeps` symmetric lexicographic sweeps over <= a few dozen cells. */ +/* The coarsest level: one thread per march, `sweeps` symmetric RED-BLACK + * sweeps (red, black, black, red — the CPU's ordering, so a nearly singular + * coarsest system settles on the same constant) from zero. */ extern "C" __global__ void mg_coarsest( - int K, int n_cells, const unsigned int* __restrict__ cells, int n, + int K, int n_cells, const unsigned int* __restrict__ cells, + int n_red, const unsigned int* __restrict__ red, + int n_black, const unsigned int* __restrict__ black, int n, const float* __restrict__ ae, const float* __restrict__ aw, const float* __restrict__ an, const float* __restrict__ as_, const float* __restrict__ ap, const float* __restrict__ b, @@ -91,10 +95,11 @@ extern "C" __global__ void mg_coarsest( int base = k * n; for (int t = 0; t < n_cells; ++t) x[base + cells[t]] = 0.0f; for (int sw = 0; sw < sweeps; ++sw) { - for (int pass = 0; pass < 2; ++pass) { - for (int q = 0; q < n_cells; ++q) { - int t = pass == 0 ? q : n_cells - 1 - q; - int g = base + cells[t]; + for (int half = 0; half < 4; ++half) { + const unsigned int* list = (half == 0 || half == 3) ? red : black; + int n_list = (half == 0 || half == 3) ? n_red : n_black; + for (int t = 0; t < n_list; ++t) { + int g = base + list[t]; float s = 0.0f; float e = ae[g]; if (e != 0.0f) s += e * x[g + 1]; float w = aw[g]; if (w != 0.0f) s += w * x[g - 1]; diff --git a/crates/specialized/rtx-cfd/tests/gpu_vcycle_bench.rs b/crates/specialized/rtx-cfd/tests/gpu_vcycle_bench.rs index f17dc1c..3ec6a3d 100644 --- a/crates/specialized/rtx-cfd/tests/gpu_vcycle_bench.rs +++ b/crates/specialized/rtx-cfd/tests/gpu_vcycle_bench.rs @@ -292,12 +292,17 @@ fn batched_device_vcycle_go_no_go() { { let lv = &mut dev[depth - 1]; let (k_i, n_cells_i, n_i, sw_i) = (k as i32, lv.n_cells as i32, lv.n as i32, 50i32); + let (n_red_i, n_black_i) = (lv.n_red as i32, lv.n_black as i32); unsafe { stream .launch_builder(&f_coarsest) .arg(&k_i) .arg(&n_cells_i) .arg(&lv.cells) + .arg(&n_red_i) + .arg(&lv.red) + .arg(&n_black_i) + .arg(&lv.black) .arg(&n_i) .arg(&lv.ae) .arg(&lv.aw)