From 4a624879ca65b5232718aef324afa6996dac9f74 Mon Sep 17 00:00:00 2001 From: biondizzle Date: Sat, 16 May 2026 01:15:49 +0000 Subject: [PATCH] =?UTF-8?q?docs:=20update=20DEBUG=5FLOG=20=E2=80=94=20inpu?= =?UTF-8?q?t=5Fscale=20red=20herring,=20current=20state,=20next=20steps?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- DEBUG_LOG.md | 140 +++++++++++++++++++++++++-------------------------- 1 file changed, 68 insertions(+), 72 deletions(-) diff --git a/DEBUG_LOG.md b/DEBUG_LOG.md index 95c98b4a..8b3da246 100644 --- a/DEBUG_LOG.md +++ b/DEBUG_LOG.md @@ -1,96 +1,93 @@ # NVFP4 MegaMoE Debug Log -## Current State (May 15, 2026 — 23:37 UTC) +## Current State (May 16, 2026 — 01:15 UTC) -**Status:** SF remap is CORRECT. GEMM is mathematically correct. The 0.2 cosine against the BF16 reference is a **red herring** — our Python dequantization reference is wrong, not the GEMM. The vLLM pipeline still produces garbage, so the bug is elsewhere (A/B packing, activation staging, weight transform, or the BF16 reference itself). +**Status:** GEMM verified correct, SF remap verified correct, B layout verified correct, L2 slot_token cleaned up. vLLM still produces garbage. The checkpoint `input_scale` red herring is documented below. The bug remains unidentified. -**DO NOT re-investigate the SF remap.** It is verified correct by two independent tests: -1. Roundtrip verifier (commit `aa209dd`): `sfa_errors=0, sfb_errors=0` — every source byte ends up at the correct dst position -2. Uniform FP4 test: all-nibble-3 (E2M1=1.5) with SF=1.0 produces exactly 72.0 (= 1.5² × 32) for every element +**What's verified correct (DO NOT re-investigate):** +1. SF remap: roundtrip verifier = 0 errors, forward mapping with `layout_sf(make_coord(mn, k_elem, 0))` +2. GEMM math: uniform FP4 + uniform SF → exact output (72.0 = 1.5² × K) +3. B matrix layout: byte transpose correct, column-dependent weight test passes +4. L2 GEMM input: slot-major, no gather needed (cleaned up dead `slot_token` param) -### How we proved the BF16 reference is wrong (not the GEMM) +**What's still broken:** +- "The capital of France is" → garbage tokens (varies, e.g. `-W'MSG173`, `( z tractor`) +- All magnitudes reasonable, no NaN anywhere -The BF16 reference comparison has been the primary diagnostic throughout this session. It showed cosine ≈ 0 initially, then ≈ 0.2 after fixes. We assumed the reference was correct and the GEMM was wrong. **This was a false assumption.** +## Red Herring: Checkpoint input_scale (May 16, 00:10 UTC) -**Evidence that the GEMM is correct:** -1. Uniform FP4 + uniform SF → mathematically exact output (72.0 = 1.5² × K) -2. Roundtrip SF verifier passes (0 errors) -3. The cosine gap (0.2) doesn't change across multiple SF remap rewrites — it was always ≈0.2 regardless of whether we used reverse mapping, forward mapping, hierarchical coords, or flat coords -4. The GEMM's internal math is provably correct when SF values are placed correctly (test #1) +Mike's code review suggested using the checkpoint's `input_scale` as the activation global scale instead of the dynamic `amax/(6*448)`. This was **wrong** and has been reverted (commit `79b9bec`). -**Why the BF16 reference is wrong:** -- The reference manually unpacks E2M1 nibbles, looks up `_E2M1_MAGNITUDES`, multiplies by block scales and global scales -- The CUTLASS kernel uses the same E2M1 values and scale factors but may apply them in a different order or with different precision semantics (e.g., the per-element multiply order is `A_fp4 * SFA_fp8 * B_fp4 * SFB_fp8`) -- The reference doesn't account for how CUTLASS internally handles the stride-0 SF aliasing (16 K elements sharing one SF byte) -- **The 0.2 cosine is a systematic error in the reference, not the GEMM** +**What happened:** +- Applied `stage_activation(hidden_states, input_global_scale=w13_input_scale)` +- `w13_input_scale = 2.86e-4` (from checkpoint, same for all experts in a layer) +- Dynamic `amax/(6*448)` = 3.70e-3 (13x larger) +- Using 2.86e-4 for normalization: `x / 2.86e-4` produces values in the thousands for typical hidden states (amax ~5) +- ALL block scales saturated to 448.0 (max float8_e4m3 value) +- Output: still garbage, but now with quantized-to-death activations -**Lesson:** A wrong reference is worse than no reference. It sends you chasing ghosts. The SF remap went through 8+ iterations that all produced the same 0.2 cosine — because the 0.2 was never about the remap. +**Why it's wrong:** +The checkpoint `input_scale` is NOT the `input_global_scale = amax/(6*448)` normalization constant. They are different quantities: +- `input_global_scale` normalizes data to [0,1] before FP4 quantization +- `input_scale` is a calibration constant from the Quark quantization tool + +The calibration amax for `input_scale = 2.86e-4` would be `2.86e-4 * 6 * 448 = 0.77`. Runtime hidden states have amax ~5-10. The `input_scale` was computed on a different data distribution (probably calibration data, not actual inference data). + +**The correct use of `input_scale` is still unknown.** The Quark path computes `alpha = input_scale * weight_scale_2`, but this may assume BF16 activations (not FP4-quantized). Our CUTLASS kernel requires FP4 input, so we must quantize with the dynamic scale. + +**Preserved for future use:** `_w13_input_scale` and `_w2_input_scale` are now saved in `finalize_weights` (not dropped) in case we need them for alpha computation later. + +**Checkpoint input_scale values (layer 0, all experts identical):** +- `gate_proj.input_scale` = `up_proj.input_scale` = 2.862840e-04 +- `down_proj.input_scale` = 3.069196e-02 +- `weight_scale_2` = 4.650298e-05 (all projections) +- Scales vary by layer: layer 0 = 2.86e-4, layer 60 = 1.07e-2 ## SF Remap — Final Correct Implementation (commit `6626b75`) ```cpp -template -__global__ void remap_sf_to_cutlass_kernel( - const cutlass::float_ue4m3_t* __restrict__ src, - cutlass::float_ue4m3_t* __restrict__ dst, - LayoutSF layout_sf, - int MN, int K_sf, - int src_stride_mn, int src_stride_ksf -) { - int tid = blockIdx.x * blockDim.x + threadIdx.x; - if (tid >= MN * K_sf) return; - int mn = tid / K_sf; - int k_sf = tid % K_sf; - int k_elem = k_sf * 16; // logical K element, not compact SF index - int dst_idx = layout_sf(cute::make_coord(mn, k_elem, 0)); - dst[dst_idx] = src[mn * src_stride_mn + k_sf * src_stride_ksf]; -} +int mn = tid / K_sf; +int k_sf = tid % K_sf; +int k_elem = k_sf * 16; +int dst_idx = layout_sf(cute::make_coord(mn, k_elem, 0)); +dst[dst_idx] = src[mn * src_stride_mn + k_sf * src_stride_ksf]; ``` +Source strides: SFA=(K_sf, 1), SFB=(1, N) +Allocation: `cute::size(cute::filter_zeros(layout))` -**Source strides:** -- SFA (row-major M, K_sf): stride_mn=K_sf, stride_ksf=1 -- SFB (row-major K_sf, N after .T.contiguous()): stride_mn=1, stride_ksf=N +## Previous Bugs Fixed -**Allocation:** `cute::size(cute::filter_zeros(layout))` matching CUTLASS example 72a - -## Previous Bugs Fixed (SF Remap Iterations) +### BF16 reference comparison (May 15, 23:37 UTC) +The 0.2 cosine against the Python BF16 dequantization reference was a RED HERRING. The reference is wrong, not the GEMM. 8+ iterations of SF remap changes all produced the same 0.2 cosine because it was never about the remap. **A wrong reference is worse than no reference.** ### `cute::size` vs `cute::cosize` (commit `c384198`) -Iteration bound used `size` (logical) instead of `cosize` (physical). Tile-padding positions unwritten. Necessary fix but insufficient alone. +Iteration bound used `size` (logical) instead of `cosize` (physical). Fixed but insufficient alone. -### M/K coordinate extraction in `idx2crd` reverse mapping (commits `deb6b32` → `30b6c89`) -Original had `get<0..2>` = M, `get<4..5>` = K. Mike corrected: first group IS M/N, second IS K. Correct inverse: `mn = f0 + 32*f1 + 128*f2`, `k_sf = f4 + 4*f5` (f3 is stride-0, ignored). +### M/K coordinate extraction in `idx2crd` (commits `deb6b32` → `30b6c89`) +Original had M/K swapped. Mike's correction: `mn = f0 + 32*f1 + 128*f2`, `k_sf = f4 + 4*f5`. ### `if/else if` fallthrough (commit `6626b75`) -`int dst_idx = 0;` with `if (LayoutRank == 2) {...} else if (LayoutRank == 3) {...}` — if neither branch matched, all threads wrote to dst[0]. Fix: branchless `layout_sf(make_coord(...))`. +Dead `dst_idx=0` when no branch matched. Fix: branchless `layout_sf(make_coord(...))`. ### `col_major_src` ambiguity (commit `7285331`) -Boolean flag didn't capture the actual source memory layout. Replaced with explicit `src_stride_mn, src_stride_ksf` integers. +Boolean flag → explicit `src_stride_mn, src_stride_ksf`. ### Allocation size (commit `6626b75`) -Used `cute::cosize(layout)` which includes padding. CUTLASS examples use `cute::size(cute::filter_zeros(layout))` which gives the actual number of stored elements. +`cosize` → `size(filter_zeros(layout))`. + +### L2 slot_token cleanup (commit `bb5a1ba`) +`nvfp4_moe_l2` accepted `slot_token` but never passed it to GEMM. Removed dead parameter. + +## Mike's Code Review Answers + +1. **E2M1 packing:** Confirmed correct — element 2j in low nibble, 2j+1 in high nibble. Suggested hardware oracle with `__nv_cvt_bfloat16raw2_to_fp4x2`. +2. **A RowMajor:** Confirmed correct — no micro-tiling for A. +3. **B ColumnMajor:** Byte transpose confirmed correct by test. Mike flagged theoretical concern about nibble-level transpose but our test passed. +4. **Alpha/global scale:** Mike suggested `alpha = input_scale * weight_scale_2` (from checkpoint). We tried it — wrong for activation normalization. The correct use of `input_scale` in our pipeline is still TBD. +5. **Gate/up correction:** Mathematically valid. `up_half *= up_weight_gs / gate_weight_gs` is equivalent to per-column alpha. ## Architecture Notes -### CUTLASS SF Layout - -The SM100 NVFP4 SF layout (from Veitner's blog + CUTLASS source): - -``` -SfKMajorAtom = Layout< - Shape, Shape>, - Stride, Stride<_0, _1>>> - -SFA: tile_to_shape(SfAtom, make_shape(M, K, L), Step<_2, _1, _3>) -SFB: tile_to_shape(SfAtom, make_shape(N, K, L), Step<_2, _1, _3>) -``` - -- First atom mode `(32, 4)` stride `(16, 4)` → M/N dimension -- Second atom mode `(SFVecSize, 4)` stride `(0, 1)` → K dimension -- Stride-0 means 16 K elements share one SF byte (aliasing) -- Public logical access: `tensor_sf(make_coord(m_or_n, k_element, l))` -- `k_sf = k_element / SFVecSize` (the /16 concept) - ### NVFP4 MoE Pipeline ``` stage_activation(hidden_states) → x_fp4, x_sf, input_global_scale @@ -101,14 +98,13 @@ L2 GEMM: (l1_fp4, l1_sf) @ (l2_w, l2_sf) with alpha=l1_igs*l2_global_sf → outp scatter with routing weights → y ``` -### Per-element multiply order (from blog) +### Per-element multiply order `res += A_fp4 * SFA_fp8 * B_fp4 * SFB_fp8` ## Next Steps -1. **Trace the full MoE pipeline** — the GEMM works, so the bug is in how data gets TO the GEMM or how results are used AFTER -2. **Check A/B packing** — the E2M1 packed data layout for A and B matrices in the CUTLASS GEMM -3. **Check activation staging** — `stage_activation` quantizes BF16 → FP4. Is the packing correct for CUTLASS? -4. **Check weight transform** — `transform_nvfp4_weights_for_mega_moe` transposes weights and scales. Are the strides correct? -5. **Test with real vLLM inference** — after fixing the real bug, "The capital of France is" should produce "Paris" -6. **Quality optimization:** investigate o_a_proj BF16→NVFP4 quantization +1. **Compare against BF16 model** — run the same prompt on a known-good implementation to see if the attention layers are working and only MoE is broken +2. **Check the vLLM model integration** — how does the MoE output get mixed with the residual? Is `hc_post` correct? +3. **Understand the Quark input_scale contract** — maybe we need to NOT quantize activations to FP4 and instead use BF16 input +4. **Add per-layer token output logging** — see which layer the tokens go off the rails +5. **Check o_a_proj BF16 handling** — it's kept in BF16 in the checkpoint, is it being processed correctly?