feat: swordfish kernels for Blackwell#1707
Merged
Merged
Conversation
…inism flag to swordfish
AlpinDale
changed the title
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Swordfish is a Blackwell-native (sm100 and sm110) GEMM kernel family descended from the same lineage as Marlin and Machete. Supported modes: symmetric GPTQ int4 (uint4b8) and int8 (uint8b128), AWQ (uint4 with zero points),
desc_act/act_order checkpoints, fused MoE, fp16 and bf16 activations, group sizes 32/64/128 and channelwise (all with the full tcgen05 prefill path; channelwise checkpoints replicate their single scale row to group 128 at weight load, a few hundred KB of duplicate scales per layer).APHRODITE_SWORDFISH_DETERMINISTIC=1forces run-stable decode reductions. Auto-selected on supported hardware; force with--linear-backend swordfish(dense) and--moe-backend swordfish(MoE). Unsupported corners (row-parallel TP with act_order, W4A8 activation quant) fall back to Marlin automatically.Benchmarks
End-to-end serving sweep against every mixed-precision backend that runs on sm100/sm110, one checkpoint per mode, 1024-token prompts, 128 output tokens, prefix caching off. Numbers are prefill tok/s (bs=1 / bs=32) and decode tok/s (bs=1 / bs=32). Backends run bf16 activations except exllama, which is fp16-only. Machete (Hopper-only) and CutlassW4A8 (activation-quant mode) cannot run on this hardware.
Per-mode tables (all backends, with footnotes)
GPTQ int4 g128 — Llama-3.1-8B
AWQ uint4 + zero points g128 — Qwen2.5-3B
GPTQ int8 g128 — Qwen2.5-3B
GPTQ int4 with desc_act — Mistral-7B-v0.2
Fused MoE, GPTQ int4 g128 — Qwen1.5-MoE-A2.7B (60 experts, top-4)
GPTQ int8 channelwise — TinyLlama-1.1B
¹ humming shipped no sm110 support (no heuristics entry, and NVML roofline queries Jetson cannot answer); this branch adds a compat shim mapping Thor onto its sm100 heuristics, which is what makes the Thor cells above possible.
² exllama reconstructs fp16 weights and runs dense GEMM above its M threshold, which briefly held the large-batch prefill cells; swordfish now ships the same strategy as its own large-M tier (all modes including zero points, act_order folded into the weight scatter) and takes those cells back with room to spare.
³ upstream TritonW4A16 bug: a qzeros shape assert trips on symmetric-GPTQ placeholder tensors.
⁴ humming's NVRTC JIT fails to compile its zero-point and int8 configurations on this stack (both boxes).
⁵ AllSpark is channelwise-int8-only — it rejects every other mode in this sweep (group 128, 4-bit, zero points, act_order, MoE), so this is its one table. It shipped arch-gated to Ampere; this branch builds it with
8.0+PTX, lifts the capability gate, and exposes it as--linear-backend allspark. Its former B200 decode lead (the M=32-48 window) is gone: the [17,48] band now dispatches to the fused atomic grid on many-SM parts, decode zeroes its output with a compute kernel instead of a memset (a per-GEMM copy-engine round trip inside captured decode graphs), and channelwise decode reads its scale row once instead of re-expanding replicated rows.⁶ exllama rejects channelwise outright ("group size must be positive"); triton's W4A16 kernel is 4-bit only.
Known residuals, all on Thor decode at bs=32 (the M 17-48 window on 20 SMs, marlin's home regime): awq -3%, moe -2%, actorder -2%, gptq4 -0.8%. Kernel-level the window sits at 0.94-1.22 of marlin after the column-quad claims; the remaining few percent needs a cooperative wide-tile decode kernel, tracked as follow-up. The same workloads win full requests on Thor regardless — prefill is 2-3x ahead in every mode.
Channelwise, kernel level
Swordfish carries channelwise by replicating the single scale row to group 128 at weight load (K/128 x N fp16, ~256 KB per 4096² layer), which buys it the full grouped machinery — tcgen05 prefill and the dense-GEMM tier. vs marlin: B200 prefill M=256-8192 1.4-4.1x, decode 0.93-1.27; Thor prefill 2.0-3.0x (int4), 1.5-2.7x (int8), decode unchanged.
Op-level three-way at 4096x4096 int8 channelwise (fp16 activations, median µs per GEMM, lower is better) shows the M-structure behind the e2e table above:
Same pattern across the other shapes: swordfish leads decode M<=16 (1.1-1.3x over both) and mid-range prefill M=256-1024 (1.4-2.6x, tcgen05); the M=17-48 band, formerly AllSpark's by 6-16%, moved to the fused atomic grid on many-SM parts and now wins or ties every cell (0.95-1.33 vs marlin, parity-plus vs AllSpark); at M>=2048 swordfish's dense tier and AllSpark converge on the same dequant-plus-cuBLAS strategy and trade places within run-to-run noise (±10%, both directions across repeat runs). Chasing this comparison on Thor exposed a stale gate: the dense tier was disabled below 100 SMs based on 4-bit measurements, but 8-bit moves twice the weight bytes through the mixed-input transform and does cross over — it now engages at M>=2048 on few-SM parts (M>=8192 for K-heavy shapes), taking Thor int8 large-M from 0.77-0.97x AllSpark to parity-or-better and 1.6-2.2x marlin.
Why
tcgen05 has no mixed-input MMA, so weight-only int4 needs a new design on this generation. NVFP4 is the main selling point of Blackwell, but there's still people out there who use GPTQ/AWQ. This is for you!