# Getting GLM-5.2 NVFP4 Post-Training off the ground

> Source: <https://patronus.ai/blog/getting-glm-5-2-nvfp4-post-training-off-the-ground>
> Published: 2026-08-21 19:37:08+00:00

# Getting GLM-5.2 NVFP4 Post-Training off the ground

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.

TL;DR: the entire setup[GLM-5.2]is a 744B-parameter MoE (~40B active). NVFP4 base, frozen (the published[nvidia/GLM-5.2-NVFP4]checkpoint), which lands as ~110 GB of weights per training GPU on 8×B200 once Transformer Engine loads it (Part I halves that) + bf16 LoRA (rank 64, MLP-only). Trainer: Megatron with tensor/expert parallelism (TP4·EP4·ETP2) on 1 node of 8×B200; for RL, rollouts are served by SGLang on a second node of 8×B200 (disaggregated, so serving never competes with the trainer for memory); orchestrated bymiles/slime. RL algorithm: GRPO, 16 samples per prompt. Reward: how far Mario travels through level 1-1, minus a time penalty, plus a flag bonus.One caveat up front:"NVFP4" isselective, not applied everywhere. Only therouted MoE expert weightsare actually in 4-bit. The attention layers, the dense early layers, the shared experts, embeddings,`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.

## The toolchain

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.

**Megatron-LMTrainer** NVIDIA's framework for training very large models.[github.com/NVIDIA/Megatron-LM ↗](https://github.com/NVIDIA/Megatron-LM)**Transformer EngineLow-precision kernels** NVIDIA's library of FP8/FP4 building blocks (quantized GEMMs, attention, LayerNorm) that actually execute the NVFP4 math on Blackwell. The "secretly-8-bit" memory bug lives here, in how it keeps a transposed copy of each weight.[github.com/NVIDIA/TransformerEngine ↗](https://github.com/NVIDIA/TransformerEngine)**SGLangRollout engine** A fast inference/serving engine. During RL it hosts the current policy and generates the rollouts (the Mario episodes). Getting it to serve an NVFP4 MoE with a LoRA overlay on the experts took a specific combination of its runner and quantization backends.[github.com/sgl-project/sglang ↗](https://github.com/sgl-project/sglang)**slimeRL framework** An open-source RL post-training framework that connects a training backend (Megatron) to a rollout engine (SGLang) over Ray, running the generate → score → learn → sync-weights loop. Notably, slime does **not** currently support LoRA: it assumes full-parameter training.[github.com/THUDM/slime ↗](https://github.com/THUDM/slime) [z.ai/blog/glm-5.2 ↗](https://z.ai/blog/glm-5.2)**milesRL framework + LoRA** The slime-derived RL framework this project builds on. Its decisive difference from slime: **miles supports LoRA**: the reason this 744B model can be adapted with a small bf16 adapter instead of full-parameter RL. But miles does **not** officially support **NVFP4**; combining 4-bit quantization with LoRA (QLoRA) is what *our* fork adds. **The stack is a chain of gaps filled, each layer supplying what the one below it lacks: slime → LoRA (miles) → NVFP4 QLoRA (our fork).** That chain is the whole reason a 744B model is trainable here at all. Most fixes in this log (dropping the second quantized copy, EP/ETP-aware loading, DAPO wiring) live in that fork's patches to miles, Megatron, and SGLang.[github.com/radixark/miles ↗](https://github.com/radixark/miles)

## Getting 744 billion parameters to train at all

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:

### The 4-bit base that was actually using 8-bit memory

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. This is a deliberate **memory-for-compute tradeoff**, and the central one of the whole project. The columnwise copy exists purely to make the backward pass *faster*: it lets the weight-gradient GEMM use the kernel's preferred layout with no runtime transformation. Dropping it buys back ~54 GB, the difference between fitting on one node and not, and in exchange the backward must dequantize and transpose on the fly every step, spending extra FLOPs to reconstruct what used to be cached. It's the QLoRA bargain in a sentence: *trade a slower backward for a model that fits.* The same tradeoff logic appears elsewhere: the DSA chunked kernels (next) stream computation to cut activation memory, and the FP4 serving path (later) pays per-forward dequant compute rather than storing a bf16 copy.

### Going deeper: why a 4-bit weight needs *two* copies (on B200)

Some notation first: let *X∈ℝ n×in* denote the input and

*W∈ℝ*the weight matrix. Blackwell's FP4 matmuls run fastest in a

out×in**TN layout**, BLAS shorthand (

**T** ransposed–

**N** ormal) for a matmul where both operands present the contraction (K) dimension as the innermost axis. In a linear layer the weight matrix is used in

*two different contractions*: the forward (

`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.

**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"](https://docs.nvidia.com/deeplearning/transformer-engine/user-guide/features/low_precision_training/nvfp4/nvfp4.html#handling-transposes) notes.

TN, NT, and what actually bindsTNis a memory-layout tag: the two GEMM inputs present the contraction dimension K as their innermost (contiguous) axis, i.e. one operand storedrow-majorand the othercolumn-major. The exact letters are library-dependent. cuBLAS names layouts by the BLAS`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, but that flexibility is about thedata, not the scales, and it doesnotdissolve the two-copy problem: NVFP4's per-16 block scales are pinned to the quantization axis, so each pass still needs its scales grouped alongits owncontraction axis. The binding constraint is the scale axis rather than the data transpose, which is why even a layout-flexible SM100 kernel still wants a separate NVFP4 quantization per contraction.

**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

[block-scaled](https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/mma_docs/tcgen05_programming.html)

`tcgen05`

**MMA** is

[K-major with a 16-wide scale vector](https://research.colfax-intl.com/cutlass-tutorial-nvfp4-blockscaled-gemm-on-nvidia-rtx-pro-blackwell-gpus-sm12x/)), 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](https://humansand.ai/blog/nvfp4-rl) (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:

### DeepSeek Sparse Attention (DSA) and its quadratic tax

GLM-5.2 uses **DeepSeek Sparse Attention (DSA)**, introduced with DeepSeek-V3.2 [(DeepSeek-AI, arXiv:2512.02556)](https://arxiv.org/abs/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](https://github.com/NVIDIA/Megatron-LM/pull/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](https://docs.nvidia.com/deeplearning/cudnn/latest/fe-oss-apis/dsa.html). 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.

### A model no kernel wanted to serve for RL training

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.

The serving gate:Three requirements had to hold simultaneously:routed MoE experts in 4-bit NVFP4,a LoRA overlay on those same experts, andan adapter re-synced after every step. Out of the box no SGLang configuration accepted all three: the runner it auto-picks on B200 rejects LoRA on MoE outright, so there is no path at all. Pinning thetritonMoE runner, serving the base as`modelopt_fp4`

, and disabling shared-experts fusion cleared the gate.

### The optimizer clobbered its own SFT checkpoint

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.

### The out-of-memory was about sequence length, not chunk size

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."

## Collapse and recovery

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, *A i=ri−mean(r1..N)*.

If every *r i* 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

**: 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.**

*degenerate attractor*### What stopped it

[DAPO](https://arxiv.org/abs/2503.14476) (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.

## Where it stands

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**

**Memory.** Transformer Engine stores two NVFP4 quantizations of every weight, one per contraction axis, so the "4-bit" base costs 110 GB per GPU. Keeping only the forward layout and dequantizing it to bf16 for the backward brings it to 56 GB and onto one node.**Attention.** The DSA reference we were pinned to materialized two O(N²) fp32 buffers, in the indexer and in the aggregation. Chunking both cut each site's peak roughly 10×, bit-identical with the flag off.**Serving.** NVFP4 experts + a LoRA overlay on them + an adapter that changes every step is a combination nothing serves by default. It needs the triton MoE runner,`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.**Warm-start.** The RL optimizer captured its fp32 master copy before the SFT adapter was loaded, so the first step wrote zeros back over it. Re-syncing the master after the load is a one-line fix and was the difference between learning and babbling.**Learning.** GRPO with an entropy bonus pinned reward to a constant: identical samples, zero group variance, zero advantage, no gradient. DAPO with the entropy term at zero held ~670 for a full 350-rollout budget.**Precision.** At the sizes we tested, 4-bit costs essentially nothing for either SFT or RL. What limits the policy is scale and single-prompt training, not the quantization.

## The code

The changes described here live in two public forks, each named for the gap it fills:

[patronus-ai/miles-nvfp4](https://github.com/patronus-ai/miles-nvfp4): the NVFP4 QLoRA layer on top of miles, including the single-copy NVFP4 path, the EP/ETP-aware adapter loader, and the DAPO wiring.[patronus-ai/megatron-lm-nvfp4](https://github.com/patronus-ai/megatron-lm-nvfp4): the trainer-side changes, chiefly the chunked DSA indexer and aggregation kernels and their gating flags.

## Related work & further reading

We're not alone in the 4-bit-RL corner. Four pointers that shaped or corroborate this work:

[The 4-Bitter Lesson](https://humansand.ai/blog/nvfp4-rl)(humans&): NVFP4 RL for weights*and*activations, framing instability as forward (policy quant error) + backward (gradient mismatch) + a handful of**sensitive weights** at the intersection. One difference in setting is worth stating plainly: theirs is**full-parameter** NVFP4 training, where every weight is updated in low precision, whereas this article is about**parameter-efficient** training, where the NVFP4 base stays frozen and only a small bf16 LoRA adapter receives gradients.[DAPO](https://arxiv.org/abs/2503.14476)(Yu et al.): the GRPO variant this recipe comes from. Its two ingredients we lean on are**dynamic sampling**, which drops groups with zero reward variance and re-samples, and** clip-higher**, which raises the upper PPO clip so low-probability tokens can still grow.[QeRL](https://arxiv.org/abs/2510.11696): NVFP4-quantized RL with LoRA (a 32B policy on a single H100, faster rollouts), and notably a finding that quantization noise can improve exploration.[Why we chose not to serve GLM-5.2 in NVFP4](https://blog.umans.ai/blog/glm-5-2-nvfp4-not-worth-serving/)(Umans AI): independently reaches the conclusion our own SGLang dig confirmed: the FP4-MoE serving path (especially with a LoRA overlay) is a rough trade today.

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:

[Pretraining Large Language Models with NVFP4](https://arxiv.org/abs/2509.25149)(NVIDIA): the definitive account of the format and how to train in it, Random Hadamard transforms, stochastic rounding, two-dimensional scaling, and the practice of holding a minority of sensitive, low-volume layers in higher precision, which is exactly the mixed-precision logic behind the "NVFP4 is selective, not blanket" caveat above.[Transformer Engine, NVFP4 "Handling transposes"](https://docs.nvidia.com/deeplearning/transformer-engine/user-guide/features/low_precision_training/nvfp4/nvfp4.html#handling-transposes): the rowwise/columnwise usage-flag scheme, straight from the library that stores both copies (and the`set_usage`

flag we scope-patch).[Colfax Research, NVFP4 Block-scaled GEMM on Blackwell](https://research.colfax-intl.com/cutlass-tutorial-nvfp4-blockscaled-gemm-on-nvidia-rtx-pro-blackwell-gpus-sm12x/): states the hard rule plainly, the`tcgen05`

MMA is**K-major** with a**16-wide scale vector** for NVFP4.[NVIDIA CUTLASS, Block-scaled GEMM](https://docs.nvidia.com/cutlass/latest/media/docs/operators/tutorials/006_block_scaled_gemm.html)and[example 72 (Blackwell NVFP4 GEMM)](https://github.com/NVIDIA/cutlass/blob/main/examples/72_blackwell_narrow_precision_gemm/72b_blackwell_nvfp4_nvfp4_gemm.cu): the reference kernels and the scale-factor layout.[DeepGEMM](https://github.com/deepseek-ai/DeepGEMM): FP8/FP4 kernels with the`NT/TN/NN/TT`

layout naming (`D = C + A @ B`

); documents SM90 (NT-only) vs SM100 (all layouts).[DeepSeek-V2](https://arxiv.org/abs/2405.04434): introduces Multi-head Latent Attention, and describes the*weight absorption*that makes it pay: at inference`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.[DeepSeek-V3.2](https://arxiv.org/abs/2512.02556)and the[V3.2-Exp tech report](https://github.com/deepseek-ai/DeepSeek-V3.2-Exp/blob/main/deepseek_v3_2.pdf): DSA itself, the lightning indexer plus fine-grained token selection that takes attention from O(L²) to O(Lk).[FlashMLA](https://github.com/deepseek-ai/FlashMLA): DeepSeek's MLA kernels, written to consume the latent directly. Worth knowing that absorption is an*inference*-oriented trade: the absorbed query is 512-wide instead of 192, so it buys memory and costs attention FLOPs, which is why serving stacks choose absorb vs non-absorb by shape and why adopting it for training is a real trade rather than a free unlock.[cuDNN DSA docs](https://docs.nvidia.com/deeplearning/cudnn/latest/fe-oss-apis/dsa.html): NVIDIA's DSA building blocks (indexer forward, top-k, sparse-attention backward, score recompute) with the SM100+ constraints on the indexer/top-k, the counterpoint to our naive Megatron`dsa.py`

.

**APPENDIX** The full ledger

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.

**Issues & resolutions across five layers (the ones not covered above)LayerIssueRoot cause → fixStatus1 · Environment** Env build (cu130/B200)pin conflicts + flaky TLS → curated venv pins, retry/mirror**solved2 · Trainingmemory** DSA indexer chunk crash2D-mask

`IndexError`

→ fixed mask indexing**solved** Slow launch (1.4 TB bridge)HF bf16 bridge re-materialized each launch → native FP4 fast-load

**open3 · Serving**

(SGLang)Slow rollouts / evalCUDA graphs + fused MoE-LoRA incompatible → concurrency (~15×)

(SGLang)

**workaround** Cold-start init errorempty LoRA path parsed as repo id → truthy guard

**solved4 · RL**

orchestrationColocate weight-sync crash

orchestration

`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 NaNSFT optimizer-state leak →

`ADAPTER_SKIP_OPTIM`

**solved** Rollout garbage / hangthinking-mode + resp-len + TP grid →

`enable_thinking=False`

, TP4·EP4·ETP2**solved5 · Warm-start**

& learningEP/ETP-unaware warm-startthe native adapter checkpoint carries no expert-parallel dimension, so every EP rank loaded block-0 → EP-aware + ETP-sliced HF load

& learning

**solved** GRPO reward collapsezero-variance groups, entropy destabilizes →

**DAPO, entropy=0solved**

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.

### Citation

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}}
```