[Perf] Choose the RDNA3 GEMM tile from the shape - #980
Open
vlluvia wants to merge 5 commits into
Open
Conversation
Replaces the hand-rolled index arithmetic in rdna3_f16_gemm with the layout-based API, and fixes a latent grid-swizzle fault found on the way. The epilogue now partitions C through a tiled_mma and make_tiled_copy_C instead of an explicit ``g_row = base + 2*si + klane`` store loop. The tiled_mma carries a permutation that reproduces the kernel's wave banding: the default stamping interleaves the repeats instead, which measured 62% slower at 3072x3072x1024. The GMEM to LDS path now goes through flat_divide plus a tiled copy with BufferCopy128b, replacing the precomputed offset tables and the flat _v8_store. The thread geometry is the same assignment the tables computed. The grid swizzle moves out of the kernel body into _group_width and _swizzle_tile_id. It derived bid_m from a group width that need not divide grid_m, so the last group addressed tiles past the end of the grid and the kernel wrote past C. Every shape in use happened to divide evenly, which is why it survived; it becomes reachable as soon as a caller asks for a tile narrower than 128x128. The block tile stays a parameter defaulting to 128x128x32, so this is a refactor: the generated ISA is byte-identical for rn and rs output across 128x128x32, 64x64x64 and 128x64x32 on gfx1100. Adds test_rdna3_wmma_atom for the gfx11 tiled copy atoms, which have to handle lanes 16-31 mirroring lanes 0-15 under the v16 operand ABI, and test_rdna3_grid_swizzle for the swizzle bijection. Co-authored-by: Cursor <cursoragent@cursor.com>
Benchmarking the migration against the unmodified kernel turned up that the swizzle bug is reachable at the default 128x128 tile, not only at narrower ones as the docstrings claimed, and that it has two symptoms rather than one. Measured on gfx1100 at 128x128x32: M of 1152, 1280 and 1664 return a wrong C, off by roughly 400x the bf16 rounding floor, while 1536 and 2560 fault the GPU. Which one you get depends on whether the address past the grid happens to be mapped, so the silent wrong answer is the common case and the fault is the lucky one. No behaviour change -- the fix itself already landed with the swizzle rewrite. This corrects the three places that understated it and adds grid_m of 9, 10 and 13 to the bijection test so the wrong-answer grids are covered alongside the two that faulted. Co-authored-by: Cursor <cursoragent@cursor.com>
vlluvia
force-pushed
the
feat/rdna3-gemm-tile-autotune
branch
from
August 7, 2026 04:40
a0a37a8 to
9550c8b
Compare
The repo formats at 120 columns; these files were wrapped nearer 100, so black wanted to rejoin several call arguments and expand the parametrize lists to one tuple per line. Formatting only, no behaviour change. Co-authored-by: Cursor <cursoragent@cursor.com>
rdna3_f16_gemm builds whatever tile it is handed and defaults to 128x128x32. That tile is right once the problem fills the grid, but it cuts only 4 workgroups at 256x256 and 16 at 512x512, so on a 96-CU part most CUs idle no matter how good the inner loop is. Choosing the tile from the shape is worth up to 3.0x there. rdna3_f16_gemm_autotune owns that decision in two layers. pick_tile is a heuristic fitted to a sweep of every feasible tile on 27 shapes; it needs no GPU and no measurement, and it is what a call resolves to with nothing configured, so the wrapper benchmarks nothing by default. Above it sits the shared autotuner: FLYDSL_AUTOTUNE=1 sweeps feasible_tiles for real, and the result can be frozen into an offline artifact. The heuristic defaults to 64x64x64 rather than the widest tile that covers the machine. Measured on gfx1100 it is fastest on 16 of the 27 shapes and holds 50-59 TFLOP/s throughout, where 128x128x32 swings between 40 and 72. Taking the widest covering tile cost up to 37% and averaged 6.5%; against the per-shape fastest tile this averages 0.6%, worst case 8.1%. Two limits worth knowing. NUM_CU is hard-coded for gfx1100, so the thresholds do not transfer to a gfx11 part with a different CU count, and shapes outside the fitted set are extrapolation -- the search exists for both cases. And _graph_bench, which captures a CUDA graph to get under the ~90us launch overhead that would otherwise swamp these kernels, still reads the multi-wave tiles a few us high below about 50us, so a tuned result for a short kernel is a hypothesis to confirm rather than a fact. feasible_tiles doubles as the search space: anything it excludes does not divide the shape, cannot fill the prefetch pipeline, or does not fit in LDS, so benchmarking it would only measure a build failure. Points the gfx11 benchmark path at the wrapper so its numbers reflect the chosen tile rather than the default. Co-authored-by: Cursor <cursoragent@cursor.com>
vlluvia
force-pushed
the
feat/rdna3-gemm-tile-autotune
branch
from
August 7, 2026 07:26
b55f0cc to
0ef0fe2
Compare
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
2 participants
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.
Why
rdna3_f16_gemmbuilds whatever tile it is handed and defaults to 128x128x32. That tile is right once the problem fills the grid, but it cuts only 4 workgroups at 256x256 and 16 at 512x512, so on a 96-CU part most CUs idle no matter how good the inner loop is. Choosing the tile from the shape is worth up to 3.0x there.How
rdna3_f16_gemm_autotuneowns the decision in two layers, leaving the kernel untouched:pick_tileis a heuristic fitted to a sweep of every feasible tile on 27 shapes. It needs no GPU and no measurement, and it is what a call resolves to with nothing configured — the wrapper benchmarks nothing by default.FLYDSL_AUTOTUNE=1sweepsfeasible_tilesfor real, and the result can be frozen into an offline artifact.feasible_tilesdoubles as the search space: anything it excludes does not divide the shape, cannot fill the prefetch pipeline, or does not fit in LDS, so benchmarking it would only measure a build failure.Why 64x64x64 as the default
Not the widest tile that covers the machine, which is the tempting rule. Measured on gfx1100, 64x64x64 is fastest on 16 of the 27 shapes and holds 50-59 TFLOP/s throughout, where 128x128x32 swings between 40 and 72. Taking the widest covering tile cost up to 37% and averaged 6.5%; against the per-shape fastest tile this heuristic averages 0.6%, worst case 8.1%.
Two limits worth knowing
NUM_CUis hard-coded for gfx1100, so the thresholds do not transfer to a gfx11 part with a different CU count, and shapes outside the fitted set are extrapolation. The search exists for both cases._graph_benchcaptures a CUDA graph to get under the ~90us launch overhead that would otherwise swamp these kernels, but it still reads the multi-wave tiles a few us high below about 50us. A tuned result for a short kernel is a hypothesis to confirm rather than a fact.Tests
test_rdna3_tile_selectionpins the heuristic against the measured shapes;test_rdna3_gemm_autotunechecks that the untuned path resolves topick_tileand that the default is reachable by the search. Both are GPU-free.test_rdna_gemmadds a device check that the wrapper agrees with the heuristic path.Full suite on gfx1100 (Radeon Pro W7900): 2216 passed, 0 failed. #979 alone, with these files removed from disk, is 1606 passed, 0 failed.
Also points the gfx11 benchmark path at the wrapper so its numbers reflect the chosen tile rather than the default.