docs: update DEBUG_LOG with M/K swap root cause

DEBUG_LOG.md

@@ -2,19 +2,36 @@
 
 ## Current State (May 15, 2026)
 
-**Status:** Root cause identified and fixed. Awaiting rebuild and test.
+**Status:** Second root cause identified — SF remap coordinate extraction has M/K swapped. Awaiting rebuild and test.
 
-**Root cause:** The SF (scale factor) remap kernel in `cutlass_nvfp4_gemm.cu` used `cute::size(layout_sf)` as the iteration bound instead of `cute::cosize(layout_sf)`. The `size` returns the logical size; `cosize` returns the physical size including tile padding. The destination buffer was allocated with `cosize` elements (correct) and zero-initialized, but the kernel only iterated over `size` elements (incorrect), leaving tile-padding positions as zero instead of their actual SF values.
+### Root Cause #1 (partially fixed): `cute::size` vs `cute::cosize` (commit `c384198`)
 
-**Why it was invisible in the all-ones test:** When all SF values are identical (uniform data), missing writes don't matter — every position should have the same value, and the ones that got written have the right one. The standalone test from the previous session used a single global scale for all blocks, producing uniform SF, which is why it showed cosine 1.0.
+The SF remap kernel used `cute::size(layout_sf)` as the iteration bound instead of `cute::cosize(layout_sf)`. This left tile-padding positions unwritten (zero). Fix: one-line change `size` → `cosize`. However, this fix alone did NOT resolve the cosine ≈ 0 problem — random data still produced garbage.
 
-**Why it broke with real data:** Different blocks have different SF values. The tile-padding positions in the CUTLASS interleaved SF layout need specific SF values, but they were left as zero. CUTLASS reads those positions during the GEMM, getting zero scales instead of the correct values, which scrambles the output direction while preserving approximate magnitude.
+### Root Cause #2 (current): M/K coordinates swapped in SF remap (commit `deb6b32`)
 
-**Fix:** One-line change in `cutlass_nvfp4_gemm.cu` line 128: `cute::size` → `cute::cosize` (commit `c384198`).
+After the cosize fix failed to resolve the issue, we ran deeper diagnostics:
+- **All-ones test (M=1, N=32, K=32):** cosine = 1.0 ✅ (uniform SF masks any coordinate bug)
+- **Random data (same dimensions):** cosine ≈ 0.2 ❌
+- **Isolated SFA and SFB remap:** both broken (cosine 0.16 and 0.21 respectively)
 
-**Original symptoms:**
-- Deterministic prompt "The capital of France is" → `-W'MSG173 ~SB…abych` instead of "Paris"
-- No NaN/Inf, magnitudes reasonable, but cosine similarity ≈ 0 between NVFP4 GEMM and BF16 reference
+The remap kernel's coordinate extraction assumed `get<0..2>` = M group and `get<4..5>` = K group. But analysis of the CUTLASS `Sm1xxBlockScaledConfig` layout reveals the opposite: the SfAtom is K-major with `Step<_2,_1>`, meaning the first atom dimension tiles along K (problem dim 1) and the second tiles along M (problem dim 0). So `get<0..2>` = K group, `get<3..5>` = M group.
+
+**Previous (wrong):**
+```cpp
+m = get<0>(flat) + get<1>(flat) * 32 + get<2>(flat) * 128;
+k_sf = get<4>(flat) + get<5>(flat) * 4;
+```
+
+**Fixed (commit `deb6b32`):**
+```cpp
+k_sf = get<0>(flat) + get<1>(flat) * 32 + get<2>(flat) * 128;
+m = get<3>(flat) + get<4>(flat) * InputSFVectorSize + get<5>(flat) * (InputSFVectorSize * 4);
+```
+
+Also added printf diagnostics in the remap kernel to print the first 10 coordinate mappings, so we can verify the extraction at runtime.
+
+**Why the M/K swap produces cosine ≈ 0 instead of just a permuted output:** The source SF data is row-major `(M, K_sf)` for SFA. If we read `src[wrong_m * K_sf + wrong_k_sf]` instead of `src[m * K_sf + k_sf]`, and the wrong indices don't correspond to valid source positions, we get completely unrelated SF values. This corrupts the per-block scaling, making the GEMM output essentially random relative to the correct answer.
 
 ## How We Found It
 
@@ -68,15 +85,18 @@ Investigated but ruled out. Both `stage_activation` and checkpoint weights use t
 [TP7] cosine=-0.010009  mse=9.0296e+00  nvfp4_amax=6.6250  ref_amax=36.500
 ```
 
-### 8. ✅ CUTLASS SF remap `size` vs `cosize` bug — ROOT CAUSE
-**Status: FIXED (commit `c384198`).** The SF remap kernel iterated over `cute::size()` (logical) instead of `cute::cosize()` (physical with tile padding). Tile-padding positions in the CUTLASS interleaved SF layout were never written and stayed zero. With uniform SF (all-ones test) the bug was invisible. With non-uniform SF (real data) it produced cosine ≈ 0.
+### 8. ✅ CUTLASS SF remap `size` vs `cosize` bug (commit `c384198`) — partial fix
+**Status: Fixed but insufficient.** Changing `size` to `cosize` was necessary (tile-padding positions were unwritten) but did NOT resolve the cosine ≈ 0 problem. The real issue was the M/K swap in coordinate extraction (hypothesis #9).
+
+### 9. ✅ SF remap M/K coordinate swap — ROOT CAUSE (commit `deb6b32`)
+**Status: FIXED, awaiting rebuild verification.** The SF remap kernel had M and K coordinates swapped in the flattened coordinate extraction. The CUTLASS `Sm1xxBlockScaledConfig` uses a K-major SfAtom with `Step<_2,_1>`, meaning `get<0..2>` maps to the K dimension and `get<3..5>` maps to the M dimension. Our code had it backwards.
 
 **How we proved it:**
-1. All-ones GEMM test (M=1, N=32, K=32): cosine = 1.0
-2. Random data GEMM test (M=1, N=32, K=32): cosine ≈ 0.2
-3. Random data sweep (multiple dimensions): cosine ≈ 0 everywhere
-4. The only difference: uniform vs non-uniform SF values → SF remap is the culprit
-5. Found `cute::size` on line 128 when comment explicitly said use `cute::cosize`
+1. `cosize` fix alone didn't resolve cosine ≈ 0
+2. All-ones test (uniform SF) still passed — coordinate bugs are invisible with uniform data
+3. Isolated SFA vs SFB: both broken (cosine 0.16, 0.21)
+4. Analyzed CUTLASS source: `Sm1xxBlockScaledBasicChunk` uses `SfKMajorAtom` where first group = K, second = M
+5. Added printf diagnostics to verify at runtime
 
 ## Key Commits
 
@@ -98,21 +118,36 @@ Investigated but ruled out. Both `stage_activation` and checkpoint weights use t
 | `363dd89` | Dimension sweep to isolate GEMM bug |
 | `60f7f60` | Ultra-minimal GEMM with all-ones (cosine=1.0!) |
 | `67dcfa8` | Random data at small dims + alpha sweep |
-| `c384198` | **FIX: SF remap uses cute::cosize() instead of cute::size()** |
+| `c384198` | Fix: SF remap uses cute::cosize() instead of cute::size() |
+| `deb6b32` | **FIX: Swap M/K in SF remap coordinate extraction + add printf diagnostics** |
 
 ## Bugs Fixed During This Debug Session
 
-### 🔴 ROOT CAUSE: SF remap `size` vs `cosize` (commit `c384198`)
+### 🔴 ROOT CAUSE: SF remap M/K coordinate swap (commit `deb6b32`)
 
-**Bug:** In `cutlass_nvfp4_gemm.cu` line 128, the SF remap kernel used `cute::size(layout_sf)` as the iteration bound instead of `cute::cosize(layout_sf)`. The `size` returns the logical element count; `cosize` returns the physical size including tile padding. The destination buffer was correctly allocated with `cosize` elements and zero-initialized, but the kernel only wrote to `size` positions, leaving tile-padding positions as zero.
+**Bug:** The SF remap kernel in `cutlass_nvfp4_gemm.cu` had M and K coordinates swapped in the flattened coordinate extraction. The code assumed `get<0..2>` = M group and `get<4..5>` = K group, but the CUTLASS `SfKMajorAtom` layout has K first and M second (K-major, with `Step<_2,_1>` tiling).
 
-**Why it was missed in the previous audit:** We changed all *allocation* sites from `size` to `cosize` (lines 179, 180, 232, 246, 287). The comment on lines 114-115 explicitly warned about this. But the *iteration bound* in the remap kernel itself (line 128) was overlooked — it was a different context (kernel launch parameter, not buffer allocation).
+**Previous (wrong):**
+```cpp
+m = get<0>(flat) + get<1>(flat) * 32 + get<2>(flat) * 128;
+k_sf = get<4>(flat) + get<5>(flat) * 4;
+```
 
-**Why the standalone test passed:** The previous standalone test used a single global scale for all blocks, producing uniform SF values. When all SF values are identical, missing writes don't matter — every position gets the same value regardless of which positions are written. The all-ones test in this session (M=1, N=32, K=32, cosine=1.0) confirmed this.
+**Fixed:**
+```cpp
+k_sf = get<0>(flat) + get<1>(flat) * 32 + get<2>(flat) * 128;
+m = get<3>(flat) + get<4>(flat) * InputSFVectorSize + get<5>(flat) * (InputSFVectorSize * 4);
+```
 
-**Fix:** `int total = cute::size(layout_sf);` → `int total = cute::cosize(layout_sf);`
+**Why the original code looked correct:** The comment said `((32, 4, n_m_tiles), (16, 4, n_k_tiles))` — M first, K second. But this is the *logical* M/K assignment, not the *physical* flattened order. The actual CUTE layout for K-major SF puts the K group first in the flattened coordinate.
 
-**Impact:** This was the root cause of all garbage output. Every GEMM call with non-uniform SF values was producing scrambled results.
+**Impact:** Every SF value was read from `src[wrong_m * K_sf + wrong_k_sf]` instead of `src[m * K_sf + k_sf]`, producing completely unrelated scale factors. The GEMM output was essentially random (cosine ≈ 0).
+
+### SF remap `size` vs `cosize` (commit `c384198`) — necessary but insufficient
+
+**Bug:** Iteration bound used `cute::size` (logical) instead of `cute::cosize` (physical). Tile-padding positions were never written.
+
+**Impact:** With uniform SF, invisible. With non-uniform SF, additional corruption on top of the M/K swap bug. Both fixes are needed.
 
 ### Weight nibble unpack reshape bug (commit `2fd55a9`)
 
@@ -167,8 +202,10 @@ scatter with routing weights → y
 
 ## Next Steps
 
-1. **Rebuild container with cosize fix** — Mike rebuilds with commit `c384198`
-2. **Run deterministic prompt** — "The capital of France is" should produce "Paris"
-3. **Run standalone random GEMM test** — should now show cosine ≈ 1.0 with random data
-4. **If output is still off:** investigate o_a_proj BF16→NVFP4 quantization (hypothesis #6)
-5. **Once working:** clean up debug prints from production code
+1. **Rebuild container with M/K swap fix** — Mike rebuilds with commit `deb6b32`
+2. **Run standalone random GEMM test** — should now show cosine ≈ 1.0 with random data
+3. **Check printf diagnostics** — verify the coordinate mapping is correct
+4. **Run deterministic prompt** — "The capital of France is" should produce "Paris"
+5. **If output is still off:** the M/K swap fix may need refinement — the `m` stride calculation (`InputSFVectorSize * 4`) may not be correct for all cases
+6. **Once working:** remove printf diagnostics from production code, clean up debug prints
+7. **Quality optimization:** investigate o_a_proj BF16→NVFP4 quantization (hypothesis #6)