[mxfp8 training] Add a fused gated-activation (SwiGLU) + MXFP8 quantization kernel - #4743
[mxfp8 training] Add a fused gated-activation (SwiGLU) + MXFP8 quantization kernel#4743wolfcomos wants to merge 1 commit into
Conversation
🔗 Helpful Links🧪 See artifacts and rendered test results at hud.pytorch.org/pr/pytorch/ao/4743
Note: Links to docs will display an error until the docs builds have been completed. This comment was automatically generated by Dr. CI and updates every 15 minutes. |
39db529 to
61eb093
Compare
|
|
||
| # Forward data and all forward scales must match the eager reference bitwise. | ||
| # Backward data may differ by one E4M3 code in a bounded fraction of elements: | ||
| # the kernel's fast sigmoid and d_silu FMA contraction have no bit-exact eager |
There was a problem hiding this comment.
do we have more context on what exactly is not matching?
There was a problem hiding this comment.
Thanks for the reviews! The sigmoid fast math caused backward not bitwise exact: https://github.com/wolfcomos/ao/blob/61eb0939b8fc4bf19ab851a023523e0c7d0c70cb/torchao/prototype/moe_training/kernels/mxfp8/cutedsl_gated_act_mxfp8.py#L308-L338.
Bitwise exact: forward qdata, forward scales, backward dUp, and backward scales (0 scale-byte diffs across 8×10⁸ elements). The entire MXFP8 quantization contributes zero delta. The mismatch (472 of 805M backward codes, 5.9e-7, all ±1 code, all in the dGate half) is from pre-quantization fp32 rounding. Sigmoid kernel uses rcp.rn(1 + ex2.approx(−x·log2e)), which differs from fp32 torch.sigmoid at 795 of 65,280 finite bf16 codes. And FMA contraction in d_silu computes dact = fma(silu, 1−s, s) in one rounding, eager does the multiply and add as two roundings.
|
I think this looks reasonable. Can we get the torchtitan PR in close-to-final form and then review both together? |
|
Thanks a lot for the reviews! Torchtitan side integration PR is now live at pytorch/torchtitan#4257. cc @vkuzo @syed-ahmed . |
… Blackwell One CuTe DSL kernel fuses the gated activation (forward and backward) with the RCEIL MXFP8 cast: gate/up are read once and the bf16 activation never round-trips through global memory. Rowwise (1x32), colwise (32x1), or both scale modes come out of a single pass, in the same blocked tcgen05 layouts as the standalone quantizers, and bit-compatible with their special-value semantics (pytorch#4725 contract: NaN/Inf amax invalidates the block with scale byte 255 and all-NaN data; byte-0 scales descale by 2^127). Ops: torchao::gated_act_mxfp8_{forward,backward} custom ops with fake impls (torch.compile fullgraph works), public wrappers, 80-case numerics suite (bitwise forward, one-code-bounded backward), and an A/B benchmark vs the unfused Triton-activation + standalone-quantizer path. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
61eb093 to
7e08b94
Compare
Summary
This PR adds one CuTe DSL kernel that fuses the swiglu with the RCEIL MXFP8 cast:
gate/upare read once, the bf16 activation tensor is never materialized, and rowwise (1x32), colwise (32x1), or both scales come outof a single pass. This kernel is bit-compatible with the standalone mxfp8 cute dsl quantizers. The consuming
composite autograd op lives in torchtitan branch pytorch/torchtitan#4257 (will refactor for upstream torchtitan PR soon), this torchao PR is kernels-only.
Performance
GB200, torch 2.14 nightly, CUDA 13.4, public DSL 4.7.0. Baseline is an already-fused-activation arm: torch.compile-fused SwiGLU (one static Triton kernel per shape) + the standalone CuTe DSL quantizers
Eager:
This host had application clocks pinned to 2062 Mhz.
End-to-end in torchtitan (DeepSeek-V3 16B, 4x GB200, EP=4, TP=1, batch 4, seq 4096, MXFP8 MoE recipe, 50 steps), three arms — the stock MXFP8 path, the fused-w13 MoE composite running the unfused activation + standalone quantizers, and the same composite backed by this PR's kernels:
Net +2.5% throughput (+3.0% vs the structurally identical unfused composite the kernel-only effect), median step time -1.8%, loss parity (max spread 3.5e-3 at step 50), and the tightest step-time tail (max step 869 ms vs 1081 stock).
The same A/B under the BF16-wgrad recipe (
mxfp8_rceil_wgrad_with_hp: the kernel emits rowwise-only casts, and the composite recomputes the activation in BF16 at the autograd boundary for the two wgrad GEMMs — a boundary with less redundant traffic for fusion to remove):The fused kernel still wins here (+1.2% throughput, -0.7% median step, kernel-only), with the margin narrowed as expected for the cheaper boundary.
API
gated_inputpacks[gate | up]— the layout a fused w13 projection produces; backward returns concatenated[dGate | dUp], so there is notorch.caton the hot path. Disabled directions come back zero-sized (fixed arity). Registered as torch custom ops with fake impls;torch.compile(fullgraph=True)works. One kernel source: direction,layout mode, and the activation (a module-level
Constexprpolicy function,_silu_pair) are compile-time parameters; shapes are runtime, so one compilation per mode serves every shape.Numerics and tests
FakeTensorMode; pointer-alignment and geometry-override guards, each with a bitwise positive case); opt-in largest-legal-shape check against the standalone quantizer. Upstream's cutedsltest_mxfp8_dsl_special_value_semanticscases pass in the same environment.Reproducing
Follow-ups
SiLU is the only activation exposed today, plan to add more activations in future PR.
the TorchAO op itself has no register_sharding or automatic fallback. The external TorchTitan composite must choose the unfused path before invoking it, so there might be kernel sharding register work need to do for torchtitan integration.