test: add opt-in shifted storage to space-time GPU prototype

tests/validation/multigrid/spacetime_gpu.cu

@@ -9,18 +9,36 @@
 #include "operators/collision_srt.cuh"
 #include "streaming/pull_double_buffer.cuh"
 
-#if DIM != 2 || NQ != 9 || COLLISION_MODEL != 0 || STORE_PRECISION != 0 || USE_DDF_SHIFTING || USE_LES || STREAMING_MODEL != 0
-#error "Space-time experiment requires unshifted FP32 D2Q9 SRT double buffers without LES"
+#if DIM != 2 || NQ != 9 || COLLISION_MODEL != 0 || STORE_PRECISION != 0 || USE_LES || STREAMING_MODEL != 0
+#error "Space-time experiment requires FP32 D2Q9 SRT double buffers without LES"
 #endif
 
 __device__ __forceinline__ void collide(float* f, float omega) {
     float rho, ux, uy, feq[NQ];
     compute_rho_u(f, rho, ux, uy);
+#if USE_DDF_SHIFTING
+    // Same perturbation-preserving SRT branch as production BlockStep.
+    const float drho = f[0]+f[1]+f[2]+f[3]+f[4]+f[5]+f[6]+f[7]+f[8];
+    const float jx = f[1]-f[2]+f[5]-f[6]+f[7]-f[8];
+    const float jy = f[3]-f[4]+f[5]-f[6]-f[7]+f[8];
+    compute_feq_from_momentum(drho, jx, jy, ux, uy, feq);
+#else
     compute_feq(rho, ux, uy, feq);
+#endif
     collide_srt_no_force(f, feq, omega);
 }
 
-__device__ __forceinline__ void health(float value, unsigned int* counts) {
+__device__ __forceinline__ double physical(double value, unsigned int q) {
+#if USE_DDF_SHIFTING
+    // Match population_storage.py's FP64 weights, not promoted FP32 d_w.
+    return value + (q == 0 ? 4.0/9.0 : (q < 5 ? 1.0/9.0 : 1.0/36.0));
+#else
+    return value;
+#endif
+}
+
+__device__ __forceinline__ void health(float stored, unsigned int q, unsigned int* counts) {
+    const double value = physical((double)stored, q);
     if (!isfinite(value)) atomicAdd(counts + 4, 1u);
     if (value <= 0.0f) atomicAdd(counts + 5, 1u);
 }
@@ -41,7 +59,7 @@ extern "C" __global__ void CoarseStart(
                        : coarse[index_f(k, q)];
     }
     collide(f, omega);
-    for (int q = 0; q < NQ; ++q) if (!covered) health(f[q], counts);
+    for (int q = 0; q < NQ; ++q) if (!covered) health(f[q], q, counts);
     stream_pull_store(k, f, out);
 }
 
@@ -56,7 +74,7 @@ extern "C" __global__ void Collide(
     if (active == nullptr || active[k]) {
         for (int q = 0; q < NQ; ++q) f[q] = in[index_f(k, q)];
         collide(f, omega);
-        if (active != nullptr) for (int q = 0; q < NQ; ++q) health(f[q], counts);
+        if (active != nullptr) for (int q = 0; q < NQ; ++q) health(f[q], q, counts);
     } else {
         for (int q = 0; q < NQ; ++q) f[q] = NAN;
     }
@@ -74,7 +92,7 @@ extern "C" __global__ void Pull(
         unsigned long j[NQ];
         compute_neighbors(x, y, j);
         stream_pull_load(k, f, in, j);
-        if (active != nullptr) for (int q = 0; q < NQ; ++q) health(f[q], counts);
+        if (active != nullptr) for (int q = 0; q < NQ; ++q) health(f[q], q, counts);
     } else {
         for (int q = 0; q < NQ; ++q) f[q] = NAN;
     }
@@ -90,6 +108,7 @@ extern "C" __global__ void Interpolate(
     unsigned int group = blockIdx.x * blockDim.x + threadIdx.x;
     if (group >= groups) return;
     const int size = sizes[group];
+    const unsigned int q = origins[group] / TOTAL_CELLS;
     for (int child = 0; child < 4; ++child) {
         double value = frozen ? (double)g0[origins[group]] : 0.0;
         if (!frozen) for (int j = 0; j < size; ++j) {
@@ -99,10 +118,10 @@ extern "C" __global__ void Interpolate(
         const float stored = __double2float_rn(value);
         pre[4u*group+child] = value;
         h[4u*group+child] = stored;
-        if (!isfinite(value)) atomicAdd(counts, 1u);
-        if (value <= 0.0) atomicAdd(counts+1, 1u);
-        if (!isfinite(stored)) atomicAdd(counts+2, 1u);
-        if (stored <= 0.0f) atomicAdd(counts+3, 1u);
+        if (!isfinite(physical(value, q))) atomicAdd(counts, 1u);
+        if (physical(value, q) <= 0.0) atomicAdd(counts+1, 1u);
+        if (!isfinite(physical((double)stored, q))) atomicAdd(counts+2, 1u);
+        if (physical((double)stored, q) <= 0.0) atomicAdd(counts+3, 1u);
     }
 }
 
@@ -142,5 +161,5 @@ extern "C" __global__ void ApplyRegisters(
     if (row >= count) return;
     const unsigned int slot = targets[row];
     predicted[slot] = __double2float_rn((double)predicted[slot] + registers[row]);
-    health(predicted[slot], counts);
+    health(predicted[slot], slot / TOTAL_CELLS, counts);
 }