Add FP8 GEMV KSplit32 scheduling for the 48-SM SM121 GPU - #32409
Tianlei Wu (tianleiwu) merged 4 commits into
Conversation
There was a problem hiding this comment.
🟡 Changes recommended
The KSplit32 dispatch may instantiate a kernel exceeding CUDA’s static shared-memory limit.
Once you've addressed the issues Copilot identified, you can request another Copilot review.
Pull request overview
Adds FP8 GEMV KSplit32 scheduling for low-SM-count SM12x GPUs.
Changes:
- Adds device- and shape-aware K-split selection.
- Adds KSplit32 kernel dispatch.
- Tests scheduling decisions.
Findings:
- Critical: Dispatch instantiates an unreachable
<32, 4>kernel requiring 64 KiB static shared memory. Launch only the reachable<32, 1>specialization directly. - Nit: Add a functional test that forces KSplit32 rather than testing only heuristic selection.
File summaries
| File | Description |
|---|---|
onnxruntime/test/contrib_ops/matmul_block_scaled_fp8_test.cc |
Tests K-split selection. |
onnxruntime/contrib_ops/cuda/math/matmul_block_scaled_fp8.cu |
Integrates selection and KSplit32 dispatch. |
onnxruntime/contrib_ops/cuda/math/matmul_block_scaled_fp8_tiling.h |
Defines scheduling heuristics. |
Review details
Suppressed comments (1)
onnxruntime/test/contrib_ops/matmul_block_scaled_fp8_test.cc:69
- This test validates only the host-side selection result. None of the existing functional cases execute the new
KSplit=32specialization on the usual SM8/SM9 test devices, so the 1024-thread launch and 32-way reduction can fail without this test noticing. Please add a functional result test with a test seam that forces KSplit 32 (covering FP16/BF16 as appropriate), rather than testing only the heuristic.
EXPECT_EQ(onnxruntime::contrib::cuda::PickFp8MmaKSplit(
c.n, c.m, c.windows, c.sm_count, c.compute_capability_major),
c.expected);
- Files reviewed: 3/3 changed files
- Comments generated: 1
- Review effort level: Balanced
💡 Configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Use FP8 GEMV output-block geometry for the SM121 thresholds and cover coherent boundaries plus measured very-wide output shapes. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
88ea12b to
da23841
Compare
There was a problem hiding this comment.
🟡 Changes recommended
The tests are not hermetic and may bypass KSplit32 when the existing kill switch is disabled.
Once you've addressed the issues Copilot identified, you can request another Copilot review.
Review details
Suppressed comments (1)
onnxruntime/test/contrib_ops/matmul_block_scaled_fp8_test.cc:134
- CUDA availability is checked only in the child, but a skipped GoogleTest process exits successfully. On a CUDA build without a usable device, the parent therefore records this test as passed instead of skipped. Check availability before spawning so the test result accurately reports that the KSplit32 kernel was not exercised.
if (!is_child_process) {
const std::string command =
"\"" + CurrentExecutablePath() +
"\" --gtest_filter=MatMulBlockQuantizedFp8WeightOpTest.GemvTensorCoreForcedKSplit32 --gtest_color=no";
ASSERT_EQ(std::system(command.c_str()), 0);
- Files reviewed: 3/3 changed files
- Comments generated: 1
- Review effort level: Balanced
Force the tensor-core GEMV path in the child process and report unavailable CUDA environments as skipped in the parent test. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Tianlei Wu (tianleiwu)
left a comment
There was a problem hiding this comment.
Reviewed the current head. The KSplit32 schedule is narrowly gated to the qualified 48-SM SM121 configuration, recursive tiles preserve the generic schedule, and the launch instantiates only the valid <32, 1> kernel. The selector boundary cases and forced FP16/BF16 test cover the new behavior, including a ragged 33-window reduction. The prior shared-memory and test-isolation concerns are addressed; I found no remaining actionable issues.
…32433) ### Description Stacked on #32409. That PR's commit is the first commit here; **only the second commit (`Pin FP8 GEMV residency ...`) belongs to this PR.** Please merge #32409 first ??? GitHub will then collapse this diff to the residency change alone. The tensor-core FP8 GEMV launches `ceil(N / 16)` blocks. A 16-warp block only fits twice per SM, so an `N` just above `32 * sm_count` spills into a second, nearly empty wave: on H200 `N = 5120` launches 1.21 waves and `ncu` measures 66% active cycles. This adds a second `__global__` entry point over the same kernel body carrying `__launch_bounds__(32 * KSplit, 3)`, and routes the shapes in that window to it. The hint caps registers at 40, which makes three blocks resident and collapses those shapes to a single wave. `cuobjdump -res-usage` confirms the new entry point at `REG:40 SHARED:9216`, so `40 * 512 * 3 = 61440` registers and `27 KiB` of shared memory per SM ??? three blocks fit. The window is narrow and the hint is a pessimization outside it, so `Fp8MmaGemvPinsResidency` in `matmul_block_scaled_fp8_tiling.h` gates on the conditions that measurement showed matter: * a grid at or below 2 blocks per SM is already one wave (`N = 1024`, 64 blocks, measured 0.96x); * a grid above 3 blocks per SM stays multi-wave either way; * the residency hint requires SM89 or newer, where FP8 tensor-core support is native. SM86 measurements below show a repeatable ~1% regression from the register cap; * 8-warp blocks (`KSplit` 8, taken from `N >= 8192`) lose 1.05-1.08x from *any* explicit `__launch_bounds__` ??? declaring it replaces nvcc's implicit bounds even when the register cap is unchanged. 32-warp blocks (the `KSplit` 32 arm from #32409) cannot host 3 blocks per SM at all; * only one row tile meets the 40-register cap. `M = 16` (two tiles) measures 0.74x and `M = 32` (four tiles) 0.24x, both from spills. The plain kernel is unchanged and stays the default for everything outside the window, so no currently-selected shape changes code path. Also adds `ORT_FP8_GEMV_KSPLIT` to override the K-split heuristic for A/B sweeps, matching the override already present in the FP4 GEMV. ### Measurements (H200, sm_90, 132 SMs, CUDA 13.0) Kernel time, FP16 activations, `block_size` 128: | N | K | before | after | speedup | |---|---|---|---|---| | 5120 | 6144 | 13.25 us | 11.04 us | 1.20x | | 6144 | 5120 | 13.18 us | 9.63 us | 1.37x | | 5120 | 17408 | 52.00 us | 36.51 us | 1.42x | On a Qwen3 27B NVFP4 weights + FP8 KV decode workload the FP8 GEMV total drops from 12150 to 11712 us per forward. End-to-end decode throughput (median of repeats, both arms built as real binaries): | config | speedup | |---|---| | batch 1, speculative | 1.030 - 1.036x | | batch 1, non-speculative | ~1.04x | | batch 4, speculative (`M = 32`, gated out) | ~1.00x | Acceptance rate is identical to four decimals in every cell. ### Measurements (GeForce RTX 5060 Ti, sm_120, 36 SMs, CUDA 13.0) Kernel medians from Nsight Systems CUDA activity traces, FP16 activations, `M = 8`, `block_size` 128. Each trace contains 100 measured launches after warmup; `before` and `after` select the plain and residency-pinned entry points from the same binary. | N | K | grid | before | after | speedup | |---|---|---|---|---|---| | 1168 | 5120 | 73 | 9.088 us | 7.936 us | 1.15x | | 1440 | 5120 | 90 | 10.928 us | 8.480 us | 1.29x | | 1440 | 6144 | 90 | 12.896 us | 10.016 us | 1.29x | | 1440 | 17408 | 90 | 34.704 us | 26.816 us | 1.29x | | 1728 | 5120 | 108 | 11.872 us | 9.264 us | 1.28x | The predicate boundaries behave as intended: `N = 1152` (exactly 2 blocks/SM), `N = 1744` (just above 3 blocks/SM), and `M = 16` all stay on the plain kernel and measure within 1% between A/B arms. A ragged in-window `N = 1169` produces bitwise-identical output (`max_abs_diff = 0`). ### Measurements (GeForce RTX 3060, sm_86, 28 SMs, CUDA 13.0) The same methodology shows that applying the hint on SM86 is a small but repeatable pessimization, which motivates the SM89 minimum: | N | K | grid | before | after | speedup | |---|---|---|---|---|---| | 912 | 5120 | 57 | 71.648 us | 72.608 us | 0.987x | | 1024 | 5120 | 64 | 72.096 us | 72.959 us | 0.988x | | 1120 | 5120 | 70 | 72.576 us | 73.183 us | 0.992x | | 1344 | 5120 | 84 | 73.760 us | 74.352 us | 0.992x | | 1024 | 17408 | 64 | 236.814 us | 239.135 us | 0.990x | Three independent `N = 1024`, `K = 5120` traces with alternating arm order reproduce a 0.988-0.991x speedup. `cuobjdump -res-usage` shows that SM86 moves from `REG:56 STACK:0` for the plain FP16 kernel to `REG:40 STACK:8` for the pinned kernel. Boundary and `M = 16` controls remain neutral, and a ragged in-window correctness case is bitwise identical. The final dispatch therefore leaves pre-SM89 devices on the plain kernel. ### Measurements (GeForce RTX 4090, sm_89, 128 SMs, CUDA 13.3) Kernel time measured with Nsight Systems CUDA traces, FP16 activations, `block_size` 128. The baseline and pinned arms are separate provider binaries built from the same merged PR head; the baseline changes only the residency dispatch to select the plain entry point. | M | N | K | Grid | Plain | Pinned | Speedup | |---:|---:|---:|---:|---:|---:|---:| | 8 | 5120 | 6144 | 320 | 18.43 us | 14.56 us | **1.27x** | | 8 | 6144 | 5120 | 384 | 16.45 us | 13.76 us | **1.20x** | | 8 | 5120 | 17408 | 320 | 101.83 us | 98.11 us | **1.04x** | Controls outside the pinned dispatch remain neutral within run-to-run variance: | Control | Plain | PR | Ratio | |---|---:|---:|---:| | `M = 8, N = 1024, K = 5120` (grid below 2 blocks/SM) | 5.76 us | 5.92 us | 0.97x | | `M = 16, N = 5120, K = 6144` (`MTiles = 2`) | 31.36 us | 30.78 us | 1.02x | Using the occurrence counts in the Qwen3.8 workload microbenchmark, the three affected shapes save about 407 us per forward. A reverse-order repeat reproduced the main gains (1.24x, 1.22x, and 1.04x respectively). ### Tests `onnxruntime_provider_test --gtest_filter='*MatMulBlockQuantizedFp8Weight*'` ??? 14/14 pass. Two new tests: * `GemvTensorCorePinnedResidencyBoundaries` ??? pure predicate test at a fixed `sm_count` of 132, asserting the `(264, 396]` column-block window, SM86 exclusion, SM89 inclusion, and the `KSplit` / `MTiles` exclusions, so the expectations do not move with the test machine. * `GemvTensorCorePinnedResidency` ??? runs the hinted kernel. Which `N` selects it depends on the device's SM count, so the shape is derived from `cudaGetDeviceProperties` (`N = 16 * (2 * sm_count + 1)`, plus a ragged `+ 5`) and cross-checked against `PickFp8MmaKSplit`. Skips on pre-SM89 devices and devices whose window lands above `N = 8192`, where the launcher drops to 8 warps and stops hinting. ### Note for #32409 `cuobjdump -res-usage` shows `MatMulBlockScaledFp8MmaGemvKernel<32, 4, ...>` at `SHARED:66560` (65 KiB). `PickFp8MmaKSplit` only returns 32 when `m <= 8`, so `MTiles` is always 1 there and that instantiation can never be launched ??? but it is still compiled and emitted. Calling `launch_mma.template operator()<32, 1>()` directly instead of going through `launch_for_ksplit<32>()` would drop it. Not changed here since it belongs to #32409.
Add a KSplit32 tensor-core GEMV specialization for
MatMulBlockQuantizedFp8Weightand select it only for qualified low-M decodeshapes on the 48-SM SM121 GB10 GPU.
The final selector preserves the generic schedule unless all of these
conditions hold:
Each FP8 GEMV output block owns 16 columns, so the thresholds use
ceil(N / 16)rather than raw N. This gives shapes with identical launchgeometry the same SM121 override: the wide boundary starts at N=16369, and the
long-reduction boundary starts at N=5105. The generic selector is unchanged,
preserving behavior on other devices.
SM120, other SM12x minors, other SM counts, M above 8, attention-sized
projections, small outputs, short reductions, and shapes immediately below
both output-block boundaries retain the generic selector.
KSplit32 dispatches directly to the only qualified launch,
<KSplit=32, MTiles=1>. This avoids instantiating unreachable<32,2>and<32,4>kernels;<32,4>would require 64 KiB of static shared memory.Requests above M=8 are recursively tiled with GB10 KSplit32 tuning disabled, so
the original request size cannot be lost during dispatch.
Kernel performance
Measured on RTX Spark (SM121, 48 SMs) as end-to-end ORT
Run()latency with anexplicit shared CUDA stream. Each final-selector cell used three alternating
process rounds, with 100 warmups and 1,000 measured calls per arm per round.
The review specifically raised the possibility that KSplit32 could regress
already-saturated, vocabulary-sized output grids. Additional three-round GB10
negative-control measurements instead showed that KSplit32 remained beneficial
as N increased through the Qwen 3.8 lm-head dimension:
Every cell used three alternating process rounds, 100 warmups, and 1,000
CUDA-event-timed ORT calls per arm. Full outputs matched with
atol=0.5andrtol=1e-3. A final-selector smoke run confirmed that the natural selectormatched forced KSplit32 at N=248320 (4.8339 ms natural versus 4.8683 ms
forced). Because the proposed regression did not reproduce and the benefit
persisted through more than 323 output waves, no arbitrary upper N cap is
applied.
Exact boundary cases were also positive across M=1, 2, 4, and 8:
N=16384), 80 K windowsN=5120), 128 K windowsThe broad initial heuristic was rejected after finding regressions at M=16/32,
on attention and small-output shapes, at 32 K windows, and around N=4096. The
final selector excludes all of those regimes. Negative controls at 1,023
output blocks (
N=16368), 79 windows, 319 output blocks (N=5104), and 127windows retain the generic schedule.
No-override runs followed the expected schedule for every selected and
excluded case. Randomized activations and varying block scales matched the
generic output in every compared full tensor.
Whole-model performance
The exact final CUDA plugin was compared with GB10 tuning disabled versus
enabled. Each value is the median paired gain from three alternating process
pairs. The code-copy workload keeps generated tokens, speculative acceptance,
tokens per target forward, and target-forward counts identical between arms.
Batch 4 is effectively neutral, with paired changes from -0.17% to +0.41%,
while every batch-1 pair improved. Median TTFT did not regress in these
matched-acceptance comparisons.