[Perf] Add wide RDNA3 GEMM tiles - #1010
Open
vlluvia wants to merge 2 commits into
Open
Conversation
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>
Extend the RDNA3 GEMM tile ladder with wide large-shape tiles that fit through an unpadded LDS layout, and pin DeepSeek-V4-Flash decode rows as explicit follow-on scope.
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.
Summary
kblockLDS layout and scheduling hints needed for the 256x256x32 path.feasible_tiles=[], with small-M/tail handling left as follow-on work.Test plan
python3 -m py_compile kernels/gemm/rdna3_f16_gemm.py kernels/gemm/rdna3_f16_gemm_autotune.py tests/unit/test_rdna3_gemm_autotune.pypython3 -m pytest tests/unit/test_rdna3_gemm_autotune.py -qbecause this environment has nopytestinstalledDependency
Depends on #980.