Getting GLM-5.2 NVFP4 Post-Training off the ground

Z.ai and NVIDIA engineers trained GLM-5.2, a 744B-parameter mixture-of-experts model quantized to 4-bit NVFP4, with reinforcement learning to play Super Mario Bros., overcoming arithmetic, distributed-systems, and training-stability defects. The setup used Megatron-LM and Transformer Engine on 8×B200 GPUs, with SGLang serving rollouts and the miles RL framework adding LoRA support, ultimately resolving reward collapse by removing a regularization term.

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