[HIP] Perf/gfx1250 per group quant - #5273
Conversation
🏷️ CI GuideRuns automatically on every PR:
Extended tests (opt-in via labels):
PR title tags & labels: |
0c78858 to
5fde45a
Compare
github-actions
Bot
changed the title
There was a problem hiding this comment.
🟡 Changes recommended
The new gfx1250 TDM staging path can issue out-of-bounds loads for extra grid blocks and the LDS staging buffer lacks required alignment for vectorized loads.
Once you've addressed the issues Copilot identified, you can request another Copilot review.
Pull request overview
Optimizes the HIP dynamic per-group quantization kernels for gfx1250 by introducing gfx1250-gated tuning choices (group traversal order, block sizing, and optional TDM staging) and a gfx1250 fast-path for bf16→fp8 conversion that uses scalef32 divide semantics.
Changes:
- Add gfx1250-gated kernel tuning: column-major group traversal for transposed-scale layouts, interleaved 32B chunking for improved memory coalescing, and launch-time block-size selection.
- Introduce an optional gfx1250 TDM (tensor) staging path that pipelines group loads through LDS for the column-major case.
- Extend store/conversion utilities to support “scale is divisor” conversion (including a gfx1250 bf16x8→fp8x8 scalef32 intrinsic path).
File summaries
| File | Description |
|---|---|
| csrc/kernels/quant_kernels.cu | Adds gfx1250-tuned group ordering, interleaved chunk gather/store, optional TDM staging, and runtime launch heuristics for block sizing/TDM enablement. |
| csrc/include/aiter_opus_plus.h | Adds gfx1250 scalef32 bf16→fp8 conversion helper, introduces scaled_cast_div, and wires a “scale_is_divisor” option through store helpers. |
Review details
- Files reviewed: 2/2 changed files
- Comments generated: 2
- Review effort level: Lite
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
| const int64_t wave_g0 = block_g0 + wave_id * kGroupsPerWave; | ||
| int64_t x0; | ||
| int32_t y0; | ||
| resolve(win_g0, x0, y0); |
| // ori_row_stride. One descriptor moves it, and the hardware picks the transaction | ||
| // sizes instead of us paying the 50% per-instruction coverage of a b128 whose | ||
| // lanes sit 32B apart. | ||
| __shared__ DTYPE_I tdm_lds[kTdmLdsWaves * kTdmRing * kTdmSlotElems]; |
[16384, 7168] bf16, group_size 128, e8m0 + transposed: 36.86 -> 26.9 us. Five independent changes: v_cvt_scalef32_pk8_fp8_bf16 for the convert (90 instructions per thread down to 4, gated on use_e8m0_scale since the hardware keeps only the scale's exponent); sinking the scale store past the conversion, which had forced a blocking s_wait_xcnt; column-major group order when the scale is transposed, making its store contiguous instead of one byte per cache line; two 32B chunks per thread instead of one 64B run, halving both lane strides; and block size from the group ordering, 64 when column-major and 256 otherwise. Bit-exact against an fp64 reference across 9 shapes; test_quant.py passes. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Wave 0 moves the block's tile global -> LDS with tensor_load_to_lds and a barrier publishes it. [16384, 7168] e8m0 + transposed: 27.0 -> 25.50 us. Only column-major order qualifies: there a block's groups are consecutive rows at one y, a plain 2D region one descriptor can move. Depth 2 from a sweep -- deeper unrolls the body once per tile and buys nothing. Gated on 6 blocks per SIMD (the switch lands between M=3072 and 4096) because staging halves the block count and costs up to 20% below that. The grid is deliberately sized by the UNSTAGED block span, leaving its back half with nothing in range. Sizing it exactly measured 15-24% slower with identical device code -- worth a dispatch-level profile before anyone "fixes" it. Bit-exact across 15 shapes and all three scale variants; test_quant.py passes. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The column-major ordering, the 32B chunk split, the 64/256 block size and TDM staging all came off gfx1250 sweeps and lean on wave32, b128 as the widest per-lane access, and that part's cache and dispatch behaviour. None was measured on gfx950, so it now keeps its pre-series shape byte for byte -- compile-checked only, no such part here. Device side gates on __gfx1250__, host side on a cached get_gpu_arch(). Block size is a template argument, so the arch picks at launch. gfx1250 unchanged (8.01/25.46 us at M = 4096/16384); correctness unchanged. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The 256-thread block divides the block count by four, which on a small tensor leaves the machine idle: at [8, 7168] that is 7 blocks on 1024 SIMDs. On the fp32-scale path it is up to 1.23x slower through T=1024 and 4-9% faster from T=2048, so it now needs 1 block per SIMD -- blocks per SIMD rather than T, the same shape as the TDM gate. This was a regression the series introduced: block size went from a constant 64 to layout-dependent 64/256 with no size condition, and the fp32-scale path measured 0.89-0.95x at T <= 16 against main. It is now 1.01-1.06x. Bit-exact across 15 shapes and all three scale variants; test_quant.py passes. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The ring never needs more slots than there are steps to stage, so at the shipped KPT of 2 one of its three slots was LDS nobody writes, 4 KiB per block. Capping it at kTdmKPT measures -3.6% at [16384, 7168] e8m0 + transposed (24.40 -> 23.53 us, 3 reps). The kernel body is byte-identical either way -- 834 instructions, zero differing lines in the disassembly -- and the output is bit-identical across 108 cases. Only the LDS request changes, 12288 -> 8192 bytes. Why that makes identical code faster is not established; it is a threshold, not a slope. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
37ded91 to
d4be180
Compare
There was a problem hiding this comment.
🟡 Changes recommended
The new gfx1250 TDM path can issue out-of-bounds async loads due to missing early-exit and unsafe host-side TDM enabling when runtime row count differs.
Once you've addressed the issues Copilot identified, you can request another Copilot review.
Review details
Suppressed comments (2)
csrc/kernels/quant_kernels.cu:397
- In the gfx1250 TDM path, the result of resolve(win_g0, ...) is ignored. Because the host intentionally oversizes the grid even when staging is enabled (see the comment in the launch code), some blocks will have win_g0 out of range; those blocks should early-exit before building a TDM descriptor and issuing async_loads, otherwise they can DMA from out-of-bounds global memory.
int64_t x0;
int32_t y0;
resolve(win_g0, x0, y0);
auto w = opus::make_tdm<TdmWindow>(
csrc/kernels/quant_kernels.cu:1252
- Host-side
use_tdmis computed usingrowsderived frominput.numel()/cols, but the kernel may overrideori_rowsfromnum_rows_ptr(device value). Ifnum_rows_ptris provided, the divisibility and tiling assumptions required by TDM (e.g., rows % kGroupsPerStagedBlock == 0) can be wrong at runtime, which risks out-of-bounds TDM loads. The simplest safe option is to disable TDM whennum_rows_ptr != nullptr.
const bool use_tdm =
dyn_gq_tuned_arch() &&
kColMajor && DYN_GQ_TDM_KPT > 1 && (rows % kGroupsPerStagedBlock) == 0 &&
(num_group / kGroupsPerStagedBlock) >= simds * kTdmMinBlocksPerSimd;
- Files reviewed: 2/2 changed files
- Comments generated: 1
- Review effort level: Lite
| static constexpr bool kStoreTakesDivisor = | ||
| use_e8m0_scale && std::is_same_v<DTYPE_O, opus::fp8_t> && | ||
| std::is_same_v<DTYPE_I, opus::bf16_t> && (thread_data_size % 8 == 0); |
`v_cvt_scalef32_pk8_fp8_bf16` has no saturation. fp8 e4m3's top two steps are
448 (0x7e) and 480, whose encoding 0x7f is taken by NaN, so the convert
round-to-nearests anything past their midpoint 464 straight into a NaN. The
software `v_med3_f32` path clamps, so the two disagreed on out-of-range input.
Two disjoint sources, both now saturated to +-448:
* An inf element made amax inf, hence an inf e8m0 scale, and turned EVERY
element of the group into NaN (128 of 128 measured). Capping amax keeps the
scale finite so the clamp can reach the inf elements. Sign is preserved,
which the med3 path did not manage -- it returned -448 for +inf.
* A finite value above 448*scale. Only reachable when the round mode floors
the scale: RoundDown / Even run up to 2x low, far outside the 3.6% headroom
up to 464. Measured with RoundDown at [4096, 7168], all-finite input:
8055 NaN out of 29,360,128 before, 0 after with max |value| 448.
NaN is deliberately left alone -- it must stay NaN, and the convert already maps
a NaN element to a NaN byte without disturbing the rest of its group. The clamp
preserves that by comparing rather than min/max-ing: both tests are false for
NaN, so it passes through untouched.
Free on the shipped default: `kScaleMayClip` is a compile-time constant, so
under RoundUp the clamp folds to `if(degenerate_group)`, a branch never taken on
finite data. Measured at [T, 7168] e8m0+transposed, three interleaved rounds:
0.4-1.9% depending on size and run, at the noise floor. VGPR unchanged
(60/52/51), and the gfx950 device assembly is byte-identical.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The three DYN_GQ_* knobs were scaffolding from the sweep. Nothing in the build ever defined them, so they were only obfuscating three constants behind the preprocessor -- with no type, no scope, and no comment at the definition. They become named constexpr values with a one-line rationale each; the sweeps they came from already live at the use sites. The comments this branch added ran to 182 lines against 289 of code. Trimmed to 122 by keeping the conclusion, the decisive number and whatever is counter-intuitive, and dropping the reasoning-out-loud, the restatements of rationale given elsewhere, and a "revisit with a profile" TODO that would only have gone stale. Pure refactor: the gfx1250 device assembly is byte-identical across both steps. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
There was a problem hiding this comment.
🔵 Needs a closer look
The PR introduces substantial arch-specific kernel control-flow and memory/launch-shape changes (including TDM staging and new conversion semantics) that warrant careful human validation beyond automated review.
Review details
- Files reviewed: 2/2 changed files
- Comments generated: 0 new
- Review effort level: Lite
Review result: 🔴 High riskI found one confirmed correctness issue in the new gfx1250 TDM path. Blocking correctness issue: runtime
|
…hold
Two dispatch-level changes, mutually exclusive by construction:
1. The transposed (column-major) path now uses blk=256 instead of blk=64
when TDM staging is not active and the grid is large enough. The wider
block covers 4x the groups and cuts the block count by 4x, which wins
4-8% from T=512 to T=4096. Below T=512 it is 2.7% negative, so the
gate sits at 1/4 block per SIMD.
2. TDM staging and the wide block cannot combine: a wider block scales the
staged tile with it (16 KiB per ring slot vs 4), and the two together
measured 22% slower at T=8192. The staging threshold is raised from 6 to
8 blocks/SIMD so staging starts at T=8192 (where it is worth ~15%)
rather than T=4096 (where it measured 10.18 vs 10.02, slightly harmful).
Measured with run_perftest, three interleaved rounds with continuous
interference monitoring (all clean), before = main merge-base:
T=512 1.02x T=1024 1.20x T=2048 1.39x
T=4096 1.36x T=8192 1.62x T=16384 1.59x
All three scale paths (e8m0+transposed, e8m0 non-transposed, fp32) are
bit-identical to the prior commit across M = 1 / 127 / 128 / 4096 / 16384 /
65536. test_quant.py exit=0 (400 items, max error 9.9e-5).
Co-Authored-By: Claude <noreply@anthropic.com>
|
There was a problem hiding this comment.
🟡 Changes recommended
The new gfx1250 TDM path can issue async loads with an out-of-range window due to the intentionally oversized grid, risking out-of-bounds global memory reads.
Once you've addressed the issues Copilot identified, you can request another Copilot review.
Review details
- Files reviewed: 2/2 changed files
- Comments generated: 2
- Review effort level: Lite
| int64_t x0; | ||
| int32_t y0; | ||
| resolve(win_g0, x0, y0); | ||
|
|
||
| auto w = opus::make_tdm<TdmWindow>( |
| // per ring slot at 256 against 4 at 64) and the two together measured 22% | ||
| // slower at T=8192. Where staging does not apply, the wide block is worth | ||
| // 4-6% from T=512 to T=4096. | ||
| static constexpr int32_t kBlkTuned = 256; | ||
| const int simds_bs = static_cast<int>(get_num_cu_func()) * 4; |
…oup-quant Co-authored-by: Cursor <cursoragent@cursor.com> # Conflicts: # csrc/kernels/quant_kernels.cu
There was a problem hiding this comment.
🔵 Needs a closer look
The gfx1250 TDM staging path can issue out-of-bounds DMA loads in intentionally over-provisioned blocks because the window-origin bounds check is computed but not enforced before creating the TDM descriptor.
Review details
Suppressed comments (1)
csrc/kernels/quant_kernels.cu:406
- In the gfx1250 TDM staging path, the grid is intentionally over-provisioned (see host-side comment) so some blocks correspond to group IDs beyond the valid range. Those blocks should return before creating the TDM descriptor / issuing async_loads; currently
resolve(win_g0, ...)is called but its boolean result is ignored, so out-of-range blocks can still form a window origin and DMA from OOB global memory.
int64_t x0;
int32_t y0;
resolve(win_g0, x0, y0);
- Files reviewed: 2/2 changed files
- Comments generated: 0 new
- Review effort level: Lite
Motivation
Technical Details
Test Plan
Test Result
gfx1250，[T, 7168] bf16，group_size=128，e8m0 + transpose_scale=True
e8m0 scale,
transpose_scale=Falsefp32 scale
Submission Checklist