fix(quant)!: keep nvfp4 global scales in fp32

[jason-fxz]

NVFP4 weight global scales (ModelOpt weight_scale_2, the reciprocal of compressed-tensors weight_global_scale) are fp32 scalars in the checkpoint, but FreeToken converted them to fp16 when it expanded them to one value per output row. A global is amax / (6 x 448), 2e-5 to 4e-4 in the checkpoints we serve, right at fp16's normal minimum (6.1e-5), so many of them became fp16 subnormals with only a few significant bits left. This PR keeps them in fp32 everywhere; the per-row layout stays as it is (gate and up, or fused dense parts, can carry different globals).

Changes

• MoE expert banks: the expert reader (models/nvfp4_banks.py), the per-row expansion (global_rows) and the Triton kernel's banks (gate_up_global, down_global) are fp32. Marlin and b12x fold the globals into their own alphas at pack time; they now start from the checkpoint values.
• Dense NVFP4 linear: the weight_global buffer and the dense readers (qwen3_5_moe, glm4_moe, gemma4, compressed-tensors in models/loader.py) are fp32.
• CPU MoE executor: the NVFP4 kernel reads the row globals as float; the fp16 decode helper it used is removed.
• Byte tables that size the banks (kernel/aot_models.py for the prebuilt fast_index_copy specs, the per-expert estimate in moe/offload_cache.py) and the ft bench bw dummy banks count 4 bytes per global.
• The GPU kernels are unchanged: they convert the loaded global to fp32 right away, and Triton specializes on the pointer dtype.

Breaking: an FTW checkpoint that stores NVFP4 global scales as fp16 is refused at load with an error that asks to reconvert it with ft checkpoint. scripts/ftw_hotfix.py reports the same problem and does not try to repair it (the fp16 rounding cannot be undone). docs/models.md and docs/ftw-hotfix.md say so.

Results

fp16 rounding of the weight globals in local checkpoints (every global tensor of the model, one value each; the error is a fixed per-row scale error of every output of that row):

checkpoint	weights	globals	in fp16 subnormal range	fp16 relative error, max
Qwen3.6-35B-A3B-NVFP4	experts	1.6e-5 - 2.4e-4	94%	1.7e-3
Qwen3.8-Flash-Next-NVFP4	experts	2.1e-5 - 4.5e-4	27%	1.2e-3
GLM-5.3-Flash-NVFP4 (RedHatAI)	experts	3.5e-5 - 3.7e-3	97%	4.9e-4
GLM-5.3-Flash-NVFP4 (LibertAI)	experts	3.5e-5 - 3.7e-3	96%	5.9e-4
Gemma-4-26B-A4B-NVFP4	experts	3.6e-5 - 5.8e-4	3%	7.8e-4
Muse-Glimmer-30B-NVFP4	dense	1.7e-5 - 1.7e-4	74%	1.5e-3
Qwen3.6-27B-NVFP4	dense	4.8e-5 - 4.2e-4	7%	5.3e-4
Qwen3.8-27B-NVFP4	dense	5.0e-5 - 4.4e-4	6%	4.6e-4
Gemma-4-31B-IT-NVFP4	dense	5.3e-5 - 2.6e-4	1%	4.2e-4

No global is below 1e-6 or rounds to zero in these checkpoints, so this is a precision fix, not a crash fix; a global below 6e-8 would have zeroed its rows.

Speed: the same kernels with fp16 vs fp32 globals (random banks, Qwen3.6-35B-A3B and GLM-4.7 shapes, triton.testing.do_bench median, two runs), RTX PRO 6000: decode (T = 1, 8) 1.000 - 1.012x, prefill (T = 4096, 16384) 0.984 - 1.013x. The global banks grow from 0.12% to 0.24% of the expert bytes.

Serving (greedy, 96 tokens): Qwen3.6-35B-A3B-NVFP4 with GPU offload and with hybrid CPU decode give coherent output that departs from main after a few lines; Qwen3.6-27B-NVFP4 and Muse-Glimmer-30B-NVFP4 give the same tokens as main.

Tests

pytest tests/moe tests/kernels tests/checkpoint tests/models tests/engine on the RTX PRO 6000: 1073 passed, 122 skipped, 27 failed; all 27 also fail on main here (19 in tests/kernels/test_ggml.py, 5 in tests/models/test_quant_config.py, test_offload.py::test_graph_capture_reuses_warm_offload_cache_before_capture, 2 in tests/models/qwen4_exp/test_weight_ckpt.py). The tests that build or check globals now use fp32, and the Qwen3.8-Flash-Next bank-size constants are updated (63.46 -> 63.63 GiB experts, 111.14 -> 111.32 GiB pinned).

[SUKUNA-AI]

Thanks for this fix. We had found the same fp16 cast downstream and carried an fp32 patch, so here are independent
measurements of its effect. They might be useful for the release notes or for anyone deciding whether the
breaking change is worth it. Setup: nvidia/Qwen3.6-35B-A3B-NVFP4 @ 1355db6a, RTX 5070 Ti 16 GB (sm_120),
FreeToken 0.1.3 plus the equivalent change, offload MoE backend, batch 1.

• Checkpoint: 28,964 of the 30,841 weight_scale_2 values (93.9%) are fp16 subnormals; only 1,434 survive the
  cast unchanged. Relative error of the cast: median 4.8e-4, p99 1.46e-3, max 1.66e-3.
• Kernels (Triton fused-MoE prefill/decode and the NVFP4 linear, real weights from 4 layers, M in
  {1, 17, 256, 2048, 4096}, against an fp64 reference): with fp32 globals, 99.97-99.99% of the bf16 outputs are
  correctly rounded; with fp16 globals, 88-94%.
• Teacher forcing (20 fixed prompts, 943 target tokens): perplexity -1.13% with fp32 globals, paired
  cluster-bootstrap 95% CI [-2.27%, -0.08%]. Other rounding-level changes we tested on the same set moved it by
  -0.37% to +0.07%.
• Functional eval (60 private cases, greedy): 59/60 with fp32 vs 57/60 with fp16. Both discordant cases favour
  fp32, but with two discordant pairs this alone is not significant.
• Caveat for anyone comparing outputs before and after: single greedy answers can flip in either direction. On
  5 smoke prompts, fp16 passed 5/5 and fp32 passed 3/5. Both "regressions" were first-token near-ties (top-2
  margins of 0.48 and 0.09 logit). An independent high-precision reference of the quantised model reproduces the
  fp32 outputs byte for byte on both prompts.
• Cost: +2 bytes per row, about 63 MB of pinned host memory for 40 x 256 experts (about 0.35% fewer
  auto-sized expert slots on a 16 GB card).