biondizzle/deepseek-v4-quant
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

docs: CORRECTED — mxf4nvf4 IS supported on sm_100a (B200)

Browse Source 
Build 17-18 'scale_vec not supported' error was because we targeted
sm_100 instead of sm_100a. The 'a' suffix enables FP4 block-scaled
instructions. No need to fall back to mxf8f6f4 with UE8M0.

Path forward: target sm_100a, use mxf4nvf4.scale_vec::4X, keep
native UE4M3 scales + block16. No scale conversion needed.
This commit is contained in:
biondizzle
2026-05-11 14:24:13 +00:00
parent cfead0012d
commit 220649c188
1 changed files with 73 additions and 47 deletions
Download Patch File Download Diff File Expand all files Collapse all files

120
README.md
View File

@@ -12,9 +12,9 @@ Full NVFP4 quantization of DeepSeek V4 Pro and vLLM serving on 8× NVIDIA B200 G
| NVFP4→BF16 Dequantization | ✅ 305 attn/shared, 91 compressor layers |
| Compressor Reconstruction | ✅ Separate kv_proj/gate_proj → fused_wkv_wgate |
| MoE Expert Serving (FusedMoE) | ✅ FLASHINFER_TRTLLM backend |
| MoE Expert Serving (MegaMoE) | 🔧 Kernel compiles, runs, but garbled (SM100 HW limit) |
| MoE Expert Serving (MegaMoE) | 🔧 Kernel compiles, runs, but garbled (wrong arch flag sm_100 vs sm_100a) |
| API Server | ✅ Running on port 8000 |
| Output Quality | 🔧 Garbled — UE8M0 scale precision loss + attention bugs |
| Output Quality | 🔧 Garbled — wrong MMA kind (used mxf8f6f4 instead of mxf4nvf4) |
## B200 Node
@@ -127,51 +127,80 @@ A native NVFP4 mega_moe kernel in our DeepGEMM fork that avoids dequantizing
expert weights to BF16 before the GEMM. The kernel keeps weights in E2M1 packed
format and uses block-scaled MMA directly.
### SM100 Hardware Limitation (CRITICAL)
### SM100a (B200) Hardware Support (CORRECTED)
**B200 (SM100) does NOT support `kind::mxf4nvf4`** (neither `scale_vec::2X` nor `4X`).
This PTX instruction requires SM103 (B300) or SM120 (GB300). On SM100, the only
FP4 block-scaled MMA is `kind::mxf8f6f4` with UE8M0 scales (block32, group_size=32).
**B200 (SM100a) DOES support `kind::mxf4nvf4`** with `scale_vec::4X` (block16, UE4M3 scales).
This is documented in PTX ISA 8.7 (CUDA 12.8+) and confirmed by NVIDIA/CUTLASS/Colfax.
| Parameter | NVFP4 Checkpoint | Kernel (SM100 Adapted) |
|-----------|-----------------|----------------------|
Build 22 produced garbled output because we incorrectly fell back to `mxf8f6f4` (UE8M0/block32),
losing 3 bits of mantissa precision per scale. The `mxf4nvf4` instruction was never actually
tried with the correct target. The real issues were:
1. **Wrong arch flag**: JIT compiled for `sm_100` instead of `sm_100a` (`a` suffix required)
2. **No NVFP4 MMA kind enum**: DeepGEMM only had BF16 + MXFP8FP4 — NVFP4 was just MXFP4 in disguise
3. **Wrong SF layout**: Block16 (scale_vec::4X) has different TMEM layout than block32 (2X)
**Fix**: Target `sm_100a`, emit `tcgen05.mma.kind::mxf4nvf4.block_scale.scale_vec::4X`,
keep E2M1 weights + UE4M3 scales + block16. No scale conversion needed. Full precision.
| Parameter | NVFP4 Checkpoint | Kernel (sm_100a, CORRECT) |
|-----------|-----------------|--------------------------|
| Weight format | E2M1 uint8 | E2M1 uint8 (unchanged) |
| Block scale format | UE4M3 (float8_e4m3fn) | UE8M0 (uint8) — **adapted for HW** |
| Block size | 16 | 32 (merged adjacent pairs, max) |
| Global scale | float32 | Folded in before UE4M3→UE8M0 |
| PTX instruction | N/A (requires SM103+) | `mxf8f6f4.block_scale` (same as MXFP4) |
| Block scale format | UE4M3 (float8_e4m3fn) | UE4M3 (native, no conversion) |
| Block size | 16 | 16 (native) |
| Global scale | float32 | Folded in before packing |
| PTX instruction | `mxf4nvf4.block_scale.scale_vec::4X` | Same |
| Instruction descriptor | `float_ue4m3_t` | Same |
**Result**: Server starts and serves, but output is **garbled**. The UE4M3→UE8M0
conversion loses 3 bits of mantissa precision per scale (8× precision loss),
which destroys output quality. The E2M1 weights are correct, but the power-of-2-only
UE8M0 scales can't faithfully represent the original UE4M3 values.
### Kernel Architecture
### Kernel Architecture (TARGET — sm_100a with mxf4nvf4)
```
sm100_fp8_nvfp4_mega_moe_impl (adapted from sm100_fp8_fp4_mega_moe_impl)
├── Same E2M1 weight packing as MXFP4
├── Same TMEM layout as MXFP4 (2X, block32)
├── Same UTCCP copy (4x32 transpose, i*4 stride)
├── mxf8f6f4.block_scale PTX instruction (UE8M0)
├── float_ue8m0_t instruction descriptor
└── UE8M0 L1 epilogue (>> 23 activation scales)
├── kGranK = 16 (NVFP4 native block size)
├── kind::mxf4nvf4.block_scale.scale_vec::4X PTX instruction
├── float_ue4m3_t instruction descriptor
├── Block16 SF layout: scale_vec::4X, 4 TMEM sub-columns per UMMA atom
├── UTCCP copy: i*8 stride (4X layout)
├── kNumSFATmemCols = SF_BLOCK_M / 32 * 4
├── kNumSFBTmemCols = SF_BLOCK_N / 32 * 4
├── kNumSFUint32 = kHidden / 64 (4 UE4M3 per int32)
├── UE4M3 L1 epilogue (float → cutlass::float_e4m3_t cast)
└── recipe = (1, 1, 16)
Python API:
├── fp8_nvfp4_mega_moe() — recipe=(1,1,32)
├── fp8_nvfp4_mega_moe() — recipe=(1,1,16)
├── transform_nvfp4_weights_for_mega_moe()
│ ├── fold_global_scale(): UE4M3 * FP32 → UE4M3
│ ├── merge_block16_to_block32(): max of adjacent pairs
│ ├── UE4M3 → float32 → UE8M0 (extract exponent byte)
│ └── pack_uint8_to_int32() + transform_sf_into_required_layout()
│ ├── NO block16→block32 merge (native block16)
│ ├── NO UE4M3→UE8M0 conversion (native UE4M3)
│ └── pack_ue4m3_to_int32() + transform_sf_into_required_layout(gran_k=16)
└── get_symm_buffer_for_nvfp4_mega_moe() — 2x SF buffer
C++ Bindings:
├── csrc/apis/mega_nvfp4.hpp
├── csrc/apis/mega_nvfp4.hpp — kGranK=16, SF stride K/16
├── csrc/jit_kernels/impls/sm100_fp8_nvfp4_mega_moe.hpp
└── csrc/apis/layout.hpp — gran_k=32 support
└── csrc/apis/layout.hpp — gran_k=16 support
```
### Current State (Build 22 — WRONG PATH)
Build 22 used `mxf8f6f4` with UE8M0/block32, which produces garbled output.
The kernel architecture below is what was deployed. It needs to be reverted to the
`mxf4nvf4` path with sm_100a targeting.
```
Build 22 (GARBLED — do not use):
├── kGranK = 32, mxf8f6f4, float_ue8m0_t
├── Same TMEM layout as MXFP4 (2X, block32)
├── UE4M3→UE8M0 scale conversion (lossy)
├── block16→block32 scale merge (lossy)
└── recipe = (1, 1, 32)
```
The earlier Build 17 code had the RIGHT instruction (`mxf4nvf4.scale_vec::4X`)
but wrong arch flag (`sm_100` instead of `sm_100a`). We need to go back to that
code and fix the arch flag.
### Container Build Pipeline
```
@@ -197,30 +226,27 @@ build_push.sh → build → login to CR → push → update docker-compose
| 12–14 | SF stride layout | Transpose to MN-major before transform |
| 15 | SymmBuffer too small (NVFP4 has 2× SF) | NVFP4-specific SymmBuffer |
| 16 | `ImportError: deep_gemm.mega.nvfp4` | Python wrapper in mega/__init__.py |
| 17 | NVCC: `scale_vec::4X` not supported on sm_100f | — |
| 18 | NVCC: `scale_vec::2X` ALSO not supported | — |
| 17 | NVCC: `scale_vec::4X` not supported on sm_100f | Wrong arch: need `sm_100a` not `sm_100` |
| 18 | NVCC: `scale_vec::2X` ALSO not supported | Same — `sm_100a` required |
| 19 | kGranK=16 still in C++ binding | → 32 |
| 20 | UE4M3→UE8M0 `uint32 >> 23` fails | Cast to int32 first |
| 22 | Server UP, but garbled output | UE4M3→UE8M0 precision loss |
| 22 | Server UP, but garbled output | Fell back to mxf8f6f4 — should use mxf4nvf4 on sm_100a |
## Path Forward
The mega_moe approach has a **hardware ceiling on B200** — the Tensor Core can't
consume UE4M3 block scales. Three options:
The `mxf4nvf4` instruction IS supported on B200 (SM100a). Build 17-18 failed because
we targeted `sm_100` instead of `sm_100a`. Build 22 garbled because we fell back to
`mxf8f6f4` unnecessarily. The correct fix:
### Option A: Fix FlashInfer FusedMoE Path (Recommended)
The `modelopt-nvfp4` branch already uses FlashInfer FP4 MoE which dequantizes
NVFP4→BF16 before GEMM. This avoids the UE4M3→UE8M0 precision loss. The garbled
output on that branch is likely from attention layer bugs (#7, #8), not the MoE.
Fix those and we should get coherent output.
### Fix the NVFP4 MegaMoE Kernel (Recommended)
1. Target `sm_100a` in DeepGEMM's JIT compiler
2. Add NVFP4 MMA kind enum (not just MXFP8FP4 with an NVFP4 hat)
3. Emit `tcgen05.mma.kind::mxf4nvf4.block_scale.scale_vec::4X`
4. Use `float_ue4m3_t` instruction descriptor
5. Block16 SF layout (scale_vec::4X) — different TMEM layout from block32 (2X)
6. Keep E2M1 weights + UE4M3 scales + block16 — no scale conversion needed
### Option B: Dequant NVFP4→BF16 in MegaMoE Shared Memory
Build a mega_moe that dequantizes in shared memory, then uses BF16 MMA.
Slower than FlashInfer but gets the mega_moe communication pattern.
### Option C: Wait for SM103+ Hardware
B300 (SM103) and GB300 (SM120) support `mxf4nvf4` natively with UE4M3 scales.
The kernel we built would work correctly on that hardware.
This gives full NVFP4 precision with zero scale format conversion.
## Running