[Perf] Choose the RDNA3 GEMM tile from the shape - #980
Conversation
b55f0cc to
0ef0fe2
Compare
|
Native gfx1100 current-main replay looks good in the scope of this PR. I replayed the final PR commit patch-id-equivalently onto current Focused results:
Kernel-only timing used HIP events around graphs containing 200 direct launches, 4 graph replays per round, 31 rounds, alternating full-ladder order. P50 selected vs pre-selection
I also checked the failing Navi workflow job. Its device test output reaches Scope boundary only, not a request to expand this PR: formal DeepSeek-V4-Flash decode rows with |
76b2399 to
588c045
Compare
|
/rerun-ci |
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>
588c045 to
6c8e7c3
Compare
|
/rerun-ci |
2 participants
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.