PERF-2 P3 bench: the coarsest level uses the CPU's red-black symmetric ordering (a nearly singular coarsest system settles on a smoother-dependent constant, which the prolongation doubled up every level: the 19 % disagreement)

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 half = 0; half < 4; ++half) {
-            for (int q = 0; q < n_cells; ++q) {
+            const unsigned int* list = (half == 0 || half == 3) ? red : black;
-                int t = pass == 0 ? q : n_cells - 1 - q;
+            int n_list = (half == 0 || half == 3) ? n_red : n_black;
-                int g = base + cells[t];
+            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];

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)