feat: add swordfish w4a16 gemm kernels for blackwell#1707
Open
AlpinDale wants to merge 10 commits into
Open
Conversation
…inism flag to swordfish
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 (all with the full tcgen05 prefill path; channelwise checkpoints route to AllSpark/Marlin, whose dequant-then-dense-GEMM strategy wins that mode at prefill).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.
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)
¹ humming has no sm110 heuristics (KeyError on Thor).
² exllama reconstructs fp16 weights and runs dense GEMM above its M threshold; on B200's dense fp16 rate that wins large-batch prefill outright (at 0.3x decode and fp16-only numerics). The same dequant-plus-dense-GEMM strategy is a candidate swordfish tier for very large M.
³ 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.
Channelwise int8 (kernel-level). No checkpoint in the sweep is channelwise, so this mode was measured at op level (4096x4096, fp16). Decode M<=16: swordfish 11.6 us vs marlin 13.8 vs AllSpark 14.1 (B200). Large M belongs to AllSpark's dequant-plus-cuBLAS path (M=8192: 185 us vs marlin 719 vs swordfish 1874, which has no channelwise prefill path). AllSpark shipped arch-gated to Ampere but its kernels run fine on Blackwell; this PR builds it with
8.0+PTX, lifts the capability gate, and routes channelwise checkpoints to it (int8) and Marlin (int4) instead of swordfish.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!