Commit Graph

5 Commits

Author SHA1 Message Date
e565ebce91 NVFP4-1.1: replace cute.math.fmin with cute.arch.fmin (correct API) 2026-05-28 03:55:54 +00:00
20d5ddfa3d NVFP4-1.1: fix indentation for @cute.jit decorators 2026-05-28 03:52:46 +00:00
f6f59d34cb NVFP4-1.1: add @cute.jit decorator to fp4_quant functions for CuTeDSL if-block support 2026-05-28 03:50:11 +00:00
6f94925491 NVFP4-1.1: fix cute.math.fmax -> cute.arch.fmax (correct CuTeDSL API) 2026-05-28 03:48:51 +00:00
80b6b79f9e NVFP4-1.1: FP4 quantization primitives for CuTeDSL kernels
- fp8_e4m3_from_float32: manual FP8 E4M3 cast (bias=7, exp 0-15 valid,
  NaN guard for exp=15/mant=7, mantissa overflow handling)
- fp8_e4m3_to_float32: dequantize FP8 E4M3 bit pattern back to Float32
- half_step_to_e2m1_idx: E2M1 step mapping (0-12 → 0-7)
- quantize_e2m1_nibble: per-element E2M1 quantize + sign + pack
- Verified 0/500 trial failures against Python reference
- Key fixes discovered during validation:
  1. FP8 E4M3 bias is 7, NOT 8
  2. Exponent range is 0-15 (exp=15/mant=7 is NaN; others valid)
  3. Subnormal formula: val = m * 2^(-9) = m/512 (NOT m/1024)
  4. Round-to-nearest-even (not round-half-up) for half_step and mantissa
  5. Mantissa overflow (round to 8) must increment exponent
2026-05-28 03:39:55 +00:00