The goal was deceptively simple to state: take GLM-5.2, a 744B-parameter mixture-of-experts model quantized to 4-bit NVFP4, attach a bf16 LoRA adapter, and train it with reinforcement learning until it could play a level of Super Mario Bros., emitting button presses, reading the terrain ahead, and running for the flag.
Realizing that specification required resolving a set of defects, which fall into three broad classes. The first is arithmetic: the 4-bit base did not fit within the memory available to it. The second is distributed-systems correctness: the warm-start adapter silently loaded the same expert block on every expert-parallel rank. The third, and the most resistant to diagnosis, concerns training stability: the reward collapse described below persisted under every standard remedy we applied, and was resolved only by removing a regularization term rather than introducing one.
lm_head, the norms, and (of course) the LoRA adapter all stay in bf16. In other words, this is a mixed-precision model, and "the NVFP4 model" is shorthand. This is deliberate on NVIDIA's part: the same "keep the sensitive, low-volume layers in higher precision" logic that the related work below leans on.
The typical post-training arc is two stages, SFT → policy-RL, and that is the sequence this article follows too.
Five pieces of open infrastructure do the work here, in two camps: a trainer that holds the weights and takes gradient steps, and a rollout engine that generates episodes, with an RL framework wiring them together.
Before any learning question could be asked, the model had to fit in memory, load successfully, and survive a forward and backward pass. A chain of infrastructure bugs stood in the way: in the environment, in the memory arithmetic, in the serving kernels, and in the distributed loader. Together they are the price of admission for QLoRA on a model this size. The standouts:
NVFP4 stores each weight in 4 bits, so after sharding across 8 GPUs the frozen base should occupy ~55 GB per GPU (744B × 4 bits ≈ 372 GB, so ~46 GB/GPU across 8 ranks, and ~55 once the per-16 block scales and the layers that stay in bf16 are counted). It consumed 110 GB per GPU, twice the expected memory. Even on a 180 GB B200, that leaves too little for activations, optimizer, and the colocated engine.
The cause: the FP4 GEMM path keeps a persistent columnwise transpose copy of the weights alongside the rowwise one, so it can perform the matrix multiplications (matmuls) of both the forward and the backward in the kernel's preferred TN layout. These aren't one tensor read two ways, they're two independent NVFP4 quantizations of the same weight, each with its own 4-bit values and its own scale set (rowwise blocks 16 along in, columnwise blocks 16 along out), stored simultaneously because the scales genuinely differ between the two blockings. Two full copies of the weight matrix, double the memory footprint, and it overflowed a single node.
The fix removes the columnwise copy, cutting the base from 110 GB to 56 GB, and reconstructs the transposed layout on the fly in a bf16 weight-gradient (dgrad) path during the backward.Quantizer carries a rowwise and a columnwise boolean, and param-init calls set_usage(rowwise=True, columnwise=torch.is_grad_enabled()), so the transpose copy is constructed simply because grad is on at build. We scope-patch the NVFP4 quantizer's set_usage to force columnwise=False during that build (not torch.no_grad(), which would strip grad_fn and break Megatron's DDP init), while still keeping the buffer for the fused-LayerNorm linears, whose backward stays on TE's native path and cannot regenerate a stripped copy. That exception costs ~1 GB against the ~54 GB recovered on the experts.
Some notation first: let X∈ℝn×in denote the input and W∈ℝout×in the weight matrix. Blackwell's FP4 matmuls run fastest in a TN layout, BLAS shorthand (Transposed–Normal)transa/transb flags, which declare whether each operand is passed transposed. "TN" is the combination where the first is and the second is not.Fprop, Y = X·Wᵀ) reads W with the input dimension as the inner dimension, while the input-gradient (Dgrad, dX = dY·W) contracts over the output dimension and wants W in the transposed layout. Transposing a packed 4-bit tensor on the fly is awkward and slow, so Transformer Engine (v2.12.0) simply pre-quantizes and stores both a rowwise and a columnwise copy of every weight matrix. Two FP4 copies, and the 4-bit base is quietly back to ~8-bit, which now occupies ~110 GB of GPU RAM.
in (a row of W), the backward's Dgrad reduces over out (a column). Blackwell's FP4 MMA (matrix-multiply-accumulate, the tensor-core instruction that multiplies two tiles and accumulates the result) requires the contraction axis innermost (K-major) with the per-16 scales grouped along it, so the forward wants an in-blocked copy and the backward an out-blocked one, two independent NVFP4 quantizations of the same weight. That is the 2× footprint; our fix keeps only the rowwise copy and reconstructs the backward in bf16.Y = X·Wᵀ reduces over in, the backward dX = dY·W reduces over out. That contracted dimension is K, the inner (matched) axis of the matmul; because the FP4 MMA requires K innermost with its 16-element scales aligned to K, and in ≠ out, the two passes need the weight in two different K-major layouts, which is why two copies exist.Why can't we just transpose the one copy? Because NVFP4 isn't plain 4-bit values, it's a block-scaled format, and the scales don't transpose. Each tensor carries its 4-bit E2M1 values plus a per-16-element-block FP8 (E4M3) scale and one FP32 global scale. An NVFP4 block is a contiguous run of 16 elements along one axis (a 1-D group, never a 2-D tile), and those blocks run along the contraction axis. The rowwise copy blocks the weight 1×16 along in (16 consecutive in values at a fixed out), while the columnwise copy blocks it 16×1 along out. A transpose changes the contraction axis, so the original scales can no longer be used: the transposed tensor requires new block groupings, and therefore new scale factors computed from the original higher-precision weights. Simply transposing the block layout along with the values does not work either, since the blocks would still be aligned to the wrong axis for the backward's contraction.
The packed FP4 values are not a free transpose view either, because they are bit-packed two per byte and swizzled for the tensor-core tile. That is why the only options are to quantize along both axes up front or, as we do, dequantize the rowwise copy to bf16 for the backward and skip the FP4 transpose entirely.
Transformer Engine documents this rowwise/columnwise scheme directly in its NVFP4 "handling transposes" notes.
transa/transb flags (hence "TN"), whereas DeepGEMM names them relative to D = C + A @ B, where its default NT (non-transposed A = row-major, transposed B = column-major) is the very same physical arrangement, e.g. fp8_gemm_nt computes D = C + A @ Bᵀ. Layout support has widened across recent GPU generations,The bottom line. Take a tiny weight [[1,2,3],[4,5,6]] quantized rowwise, one scale per row: 1 is stored under scale 3, 4 under scale 6. Dgrad reads columns, so it needs a single scale for the block {1,4}, but those two elements were quantized under different row-scales, so the existing scales cannot produce the correct column scale max(1,4)=4. Recovering it means dequantizing each element back to bf16 and re-quantizing, which needs the full-precision values already discarded by FP4. The distinction is worth keeping straight: that Dgrad contracts over out is a mathematical requirement; the FP4 kernel then demanding per-16-block scales along that contraction axis is a hardware constraint (Blackwell's tcgen05 block-scaled MMA is K-major with a 16-wide scale vector), and it binds only if you run Dgrad on the FP4 datapath. Both real fixes accept that rule rather than dodge it: Transformer Engine satisfies it by storing the second, out-major copy (the 110 GB), while we sidestep the datapath, dequantizing to bf16 and running Dgrad on the unrestricted bf16 tensor cores. What you cannot do is talk the FP4 MMA into reading the rowwise scales for the other contraction.
For a frozen QLoRA base the columnwise copy is entirely overhead: the base weights never update, so we can construct the transpose in the backward instead of caching it. That aligns with what the 4-Bitter Lesson (humans&) arrives at from the other direction: their Transformer Engine work "avoids storing additional quantized tensor copies," cutting training peak memory ~70%, paired with a dequantized backward (differentiating DQ(Q(w)), the exact quantized weights, in bf16) so the backward matches the forward's rounding decisions and gradients stay stable. There is a second and more structural reason the bf16 detour works at all: bf16 carries no per-block scales to pin to an axis, and Blackwell accepts all operand layouts rather than only TN, so the K-major constraint that forces two FP4 copies simply does not arise once the weight is dequantized. The backward gets its transposed view for free. Ours is the QLoRA case of the same idea: drop the columnwise copy, dequantize in the backward, and the "4-bit" model finally fits on a single node.
Each fix was validated first on 5-layer and 10-layer "truncated" testbeds, cheap enough to run hundreds of RL rollouts in hours, before being spent on the full 744B, where every launch costs ninety minutes just to load.
And the question arises: does 4-bit cost anything in terms of degrading learning quality? For the SFT, essentially nothing. Running the same behavior-cloning SFT on a controlled small model at both precisions, the loss curves are all but indistinguishable:
GLM-5.2 uses DeepSeek Sparse Attention (DSA), introduced with DeepSeek-V3.2 (DeepSeek-AI, arXiv:2512.02556), which pairs a "lightning indexer" with fine-grained token selection so each query interacts with only a subset of past tokens. The Megatron implementation we were pinned to (the naive reference in dsa.py as it stood in June 2026, since rewritten by PR #5099 to Megatron-LM) hid two O(N²) fp32 scratch buffers: one in the indexer, one in the attention aggregation (the larger hog). At the sequence lengths a Mario episode produces, those buffers alone blew the memory budget. The fix was a pair of gated chunked kernels (one for the indexer, one for the attention aggregation) that stream the computation in blocks, cutting each site's peak by roughly 10× while remaining bit-identical with the flag off. To be clear, this is a property of the naive reference we were on, not of DSA in general: NVIDIA does ship DSA kernels. Chunking simply made the unmodified Megatron dsa.py fit until a fused path (FlashMLA, a tilelang SparseMLA plugin, or our own gather-based sparse-flash port) could be adopted. Megatron-LM has since merged exactly that: PR #5099 (merged 2026-07-08) wires cuDNN's DSA kernels and the FlashMLA forward into core and, with them, adds tensor- and context-parallel support for DSA.
But doesn't FlashAttention already solve this? Only for a different O(N²): FlashAttention tiles the softmax so the dense N×N score matrix is never materialized, but our two hotspots live in DSA's extra machinery, which its kernel doesn't cover. The lightning indexer scores every query×key pair to decide the top-k selection (upstream of attention), and the aggregation combines over the selected tokens; the naive dsa.py built both as explicit fp32 O(N²) scratch. DSA exists to make attention O(N·k), yet its naive implementation reintroduced O(N²) memory. Our chunking is really the FlashAttention idea (stream in blocks, don't materialize) applied to the stages the flash kernel leaves out.
RL needs the policy served, not just trained. Every step has to generate episodes before there is anything to score, which stacked three requirements that no serving stack expected to see together: routed MoE experts in 4-bit NVFP4, a LoRA overlay on those same experts, and an adapter that changes after every optimizer step. Out of the box, no SGLang configuration accepted that combination, so for a while nothing would serve the model we had just finished training.
What cleared the gate was a combination of settings rather than a new kernel. The MoE runner SGLang auto-picks on B200 refuses outright, reporting that LoRA on MoE is unsupported, so the engine has to be pinned to the triton runner, which accepts it. The base is then served as modelopt_fp4 with shared-experts fusion disabled, because otherwise the bf16 shared expert folds into the packed FP4 buffer. A smaller papercut sat behind all of it: on a cold start with no adapter, an empty LoRA path was parsed as a HuggingFace repo id and the engine failed trying to download it. A truthy guard fixed that one.
Serving works now, but two of the optimizations you would normally reach for are off the table. CUDA graphs, usually the single largest win for decode, is incompatible with the FP4 MoE-LoRA path, and so is the fused MoE-LoRA kernel that makes a LoRA overlay on MoE cheap in the first place. What is left is request concurrency, worth roughly 15×, and that is the whole budget. For an inference deployment this would be an acceptable trade. In RL it lands straight on wall-clock per iteration, because rollout generation sits on the critical path of every single step. Two further consequences: we run the engine disaggregated on a second node rather than colocated with the trainer, after a memory-saver crash during the first weight sync, and we cap the response length, which is what the OOM below turns out to be about.
That second incompatibility is worth one more sentence, because it is not only a throughput fact. With the fused kernel unavailable, the engine has to apply the adapter through an explicit non-fused NVFP4 MoE-LoRA path. So losing the fused kernel does not just cost throughput, it decides which code path applies the LoRA delta. That turned out to matter: because the delta is computed in the open rather than inside a fused kernel, it was instrumentable, and it is where we eventually caught the adapter going to zero. At the first weight sync the MoE-LoRA delta collapsed by roughly 400×, driven entirely by the per-expert LoRA-B norm (7.15 down to 0.061) while the shared LoRA-A norm barely moved (192 to 191). That reading is what turned the warm-start bug below from a guess into a diagnosis.
modelopt_fp4, and disabling shared-experts fusion cleared the gate.
The most insidious bug produced a symptom that looked like everything else: rollout 0 scored well, then rollout 1 collapsed to a floor value and the model babbled. The cause was an ordering hazard. A fresh RL optimizer captures its fp32 master copy of the weights at construction time, before the warm-start adapter is loaded from disk. The master therefore held the LoRA-B initialization (all zeros). The first optimizer.step() faithfully copied that stale master back into the model, zeroing the routed-expert adapter we had just loaded. Every subsequent weight-sync then shipped all-zero adapter weights to the inference engine, which ran the pure base model and produced degenerate repetition.
The fix is one line, run after the adapter loads: optimizer.reload_model_params(), re-syncing the master from the freshly-loaded weights. It is the reason anything downstream works at all.
optimizer.step() copies the stale all-zeros master back over the loaded adapter, so from rollout 1 the engine runs the pure base model and reward falls to a ~−170 floor it never leaves. With reload_model_params(), the master is re-synced from the freshly loaded adapter.The trainer kept running out of memory, but not deterministically. It would clear five rollouts, then die on the sixth. The instinct was to shrink the DSA sparse-attention chunk size; it didn't help. The real driver was the sampled episode length: a rollout that happened to draw a long Mario run produced a long token sequence, and the DSA indexer's activation memory scales with the square of that length. Whichever rollout first drew a long episode blew the budget. Capping the response length bounded the peak, and the runs held.
Part I in one line
A 4-bit base that used 8-bit memory, sparse attention with a quadratic tax, a model no kernel would serve, a master-copy race, an EP/ETP-unaware loader, and a length-driven OOM: cleared, one by one, between "the model exists" and "the model takes a gradient step."
With the infrastructure solid, the 744B trained end-to-end. The real question surfaced: does it learn? For a long time, the answer was a specific and maddening kind of no.
Every run followed the same pattern. Reward would climb (sometimes spectacularly, past 1000, well above the behavior-cloned starting point) and then, without warning, fall off a cliff to a constant value and stay there. It did not oscillate. It did not gradually degrade. It simply snapped to a single number, exactly, rollout after rollout.
That the reward became exactly constant (261.5, then later 362.0, to four significant figures) was the whole clue. A constant group reward means all sixteen samples in a GRPO group are byte-identical. And GRPO computes each sample's advantage as its reward minus the group mean, Ai=ri−mean(r1..N).
If every ri is equal, every advantage is zero, and the policy-gradient term vanishes. The only gradients left are the weak regularizers, and they simply nudged the policy from one deterministic mode into a slightly different one. The model had fallen into a degenerate attractor: confident, deterministic, and unable to generate the diversity it needed to escape. The confirming fingerprint was the rollout log-probability crashing to ≈0 (probability ≈ 1 on every token) precisely at the collapse.
DAPO (Yu et al., arXiv:2503.14476) is the GRPO variant built for this failure mode: its exploration comes from clip-higher, raising the upper PPO clip so low-probability tokens can still grow. We had been adding an entropy term on top of it, at coefficient 0.1, which is not DAPO. Removing that term, while keeping dynamic sampling to drop the zero-variance groups, gives the green curve in Figure 7: ~670 reward held for the entire 350-rollout budget, with no collapse. The rollout log-probability also stayed comfortably below zero throughout, which is the concrete thing to watch: at each step the policy still assigned real probability to more than one continuation, instead of putting probability ≈ 1 on a single token.
The counter-intuitive part is the direction of the effect. The entropy bonus is meant to prevent exactly this collapse, and once groups went degenerate it became the dominant gradient and pushed the policy deeper into the frozen mode. Removing it was the decisive change.
That recipe now trains 744B with LoRA, with all of Part I's fixes underneath it.
Precision parity holds in RL too, not just SFT. Run the identical policy recipe on a bf16 and an NVFP4 5-layer base and the reward curves track each other (Figure 9): the 4-bit base is essentially free for the RL as well, so the flatness at 744B is about scale, not precision.
The through-line is what matters. A model too large to fit on its hardware, loaded wrong across sixteen GPUs, that no serving stack would run as configured, and trapped in a degenerate RL attractor now trains end-to-end and holds its policy. None of it required a new algorithm. It required getting the arithmetic right, the memory layout right, and the loader right, and then having the discipline to remove a term rather than add one.
In summary
modelopt_fp4, and shared-experts fusion off, and it gives up CUDA graphs and the fused MoE-LoRA kernel, leaving request concurrency as the only speed lever.The changes described here live in two public forks, each named for the gap it fills:
We're not alone in the 4-bit-RL corner. Four pointers that shaped or corroborate this work:
And the primary sources behind the two-copy / transpose story in Part I, worth reading if you want the format and kernel details first-hand:
set_usage flag we scope-patch).tcgen05 MMA is K-major with a 16-wide scale vector for NVFP4.NT/TN/NN/TT layout naming (D = C + A @ B); documents SM90 (NT-only) vs SM100 (all layouts).W^UK folds into the query up-projection and W^UV into the output projection, so only the compressed latent is ever cached. Absorption stays valid because MLA decouples RoPE onto a separate small slice that never touches the latent. This is the reordering our fused DSA path requires (absorbed_mla=True), and the reason context parallelism all-gathers a 576-channel latent rather than expanded per-head K/V.dsa.py.The issues that had to be cleared but did not earn a section of their own, by layer of the stack. Three remain open, all of them optimizations or nice-to-haves rather than blockers.
| Layer | Issue | Root cause → fix | Status |
|---|---|---|---|
| 1 · Environment | Env build (cu130/B200) | pin conflicts + flaky TLS → curated venv pins, retry/mirror | solved |
| 2 · Training memory | DSA indexer chunk crash | 2D-mask IndexError → fixed mask indexing | solved |
| Slow launch (1.4 TB bridge) | HF bf16 bridge re-materialized each launch → native FP4 fast-load | open | |
| 3 · Serving (SGLang) | Slow rollouts / eval | CUDA graphs + fused MoE-LoRA incompatible → concurrency (~15×) | workaround |
| Cold-start init error | empty LoRA path parsed as repo id → truthy guard | solved | |
| 4 · RL orchestration | Colocate weight-sync crash | torch_memory_saver / illegal-address at r0→1 → run non-colocate | workaround |
| 744B RL OOM (3 modes) | trainer-backward / adapter-receive / frag → chunk + SGLANG_MEM + GC | solved | |
| Stage-2 RL NaN | SFT optimizer-state leak → ADAPTER_SKIP_OPTIM | solved | |
| Rollout garbage / hang | thinking-mode + resp-len + TP grid → enable_thinking=False, TP4·EP4·ETP2 | solved | |
| 5 · Warm-start & learning | EP/ETP-unaware warm-start | the native adapter checkpoint carries no expert-parallel dimension, so every EP rank loaded block-0 → EP-aware + ETP-sliced HF load | solved |
| GRPO reward collapse | zero-variance groups, entropy destabilizes → DAPO, entropy=0 | solved |
Open items (all non-blocking): native FP4 fast-load to skip the 90-minute bridge; a proper EP-aware native save format; and restoring colocate mode once the memory-saver crash is fixed.
Cited as:
Fujinuma, Yoshinari, Zhe Li, Varun Prashant Gangal, and Mariya I. Vasileva. "Getting GLM-5.2 NVFP4 Post-Training off the ground." Patronus AI Engineering, July 2026.
Or in BibTeX:
@misc{fujinuma2026glm52nvfp4,
title = {Getting GLM-5.2 NVFP4 Post-Training off the ground},
author = {Fujinuma, Yoshinari and Li, Zhe and Gangal, Varun Prashant and Vasileva, Mariya I.},
year = {2026},
month = {July}
}
Patronus AI, July 2026.
← Back to Patronus Insights