[smoke] Retarget the generic mode block size tests and move them to smoke-dev - #2480
Conversation
…ead's warp These five tests pin ConstWGSize to 257, the block size a generic kernel got when the compiler widened it by one thread for the main thread. ROCm/llvm-project#4196 takes the upstream layout, which reserves a whole warp instead, so the same kernels now report 288 on a wave32 target and 320 on wave64. Accept either rather than a single number, since the value depends on the wavefront size of the GPU the test runs on. Co-authored-by: Cursor <cursoragent@cursor.com>
There was a problem hiding this comment.
Pull request overview
Updates AOMP smoke-test FileCheck expectations to accommodate ROCm/llvm-project#4196’s upstream generic mode block layout change, where generic kernels reserve a full warp (not a single thread) for the main thread—changing reported ConstWGSize depending on wavefront size.
Changes:
- Relax
ConstWGSizechecks from a single pinned value (257) to accept288(wave32) or320(wave64) for the affected generic kernels. - Apply the updated check across five smoke tests that were previously tied to the old block-size accounting.
Reviewed changes
Copilot reviewed 5 out of 5 changed files in this pull request and generated no comments.
Show a summary per file
| File | Description |
|---|---|
| test/smoke/no-loop-split-3/no_loop_split_3.c | Updates the ConstWGSize expectation for the SGN:3 kernel to accept wave32/wave64 outcomes. |
| test/smoke/no-loop-split-2/no_loop_split_2.c | Same ConstWGSize expectation update for the corresponding SGN:3 kernel. |
| test/smoke/big-jump-loop-split-4/big_jump_loop_split_4.c | Adjusts the SGN:3 kernel ConstWGSize check to tolerate both wave sizes. |
| test/smoke/big-jump-loop-split-3/big_jump_loop_split_3.c | Relaxes the SGN:3 kernel ConstWGSize check to `288 |
| test/smoke/big-jump-loop-split-2/big_jump_loop_split_2.c | Relaxes the SGN:3 kernel ConstWGSize check to `288 |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
…e tests These two tests pin ConstWGSize to 257 for their generic kernels, and to 961/993 and 129 for the kernels carrying an explicit -fopenmp-gpu-threads-per-team. ROCm/llvm-project#4196 reserves a whole warp for the main thread rather than a single thread, so the same kernels now report 320, 1024 and 192 on a wave64 target. Accept either wave size wherever the value depends on it. The 1024-thread kernel in option1 clamps to 1024 on both wave sizes, so its two wave-specific CHECK lines collapse into one and the Makefile no longer has to select a check prefix. Co-authored-by: Cursor <cursoragent@cursor.com>
|
@ronlieb could you merge this when you get a chance, if you haven't already? ROCm/llvm-project#4196 fails CI without it. The All seven pass locally on an MI210 against a build of #4196: |
dhruvachak
left a comment
dhruvachak
left a comment
There was a problem hiding this comment.
The changes themselves look good to me.
One concern is that now the older compilers will fail these tests. Do we have proper tagging of these tests so that older compilers see the older versions of these tests? Otherwise, these tests may have to be moved to smoke-dev. @ronlieb What do you think?
i think moving to *-dev is next step. |
smoke and smoke-limbo are not versioned: they have to pass against 7.14.x, 10.0, 10.1 and whatever else is current, so a test whose expected value moved with the compiler cannot stay there. The two preceding commits retargeted these seven at the warp-reserving block size ROCm/llvm-project#4196 produces, which makes them fail on every compiler that predates it. Move them to smoke-dev, which does not block, following no-loop-4 and no-loop-7 in aa1aa88. They come back in twelve weeks, once the compilers in the field produce the new geometry. Co-authored-by: Cursor <cursoragent@cursor.com>
|
Good point on versioning — done in The five from I left |
lfmeadow
changed the title
Companion to ROCm/llvm-project#4196, which takes upstream's generic mode block
layout: a pure generic kernel's block is widened by a whole warp for the main
thread rather than by the single thread we added. Seven tests pin the old
geometry and fail the moment that change lands —
no-loop-split-2,no-loop-split-3andbig-jump-loop-split-2/3/4insmoke, plusworkgroup_size_option1andworkgroup_size_option2insmoke-limbo. Allseven still print
Success; it is only theConstWGSizein their FileChecklines that is stale.
The first two commits retarget them. A 256-thread generic kernel goes from 257
to 288 on wave32 and 320 on wave64, so those checks accept either rather than a
single number;
workgroup_size_option1's 1024-thread kernel goes from 961/993to 1024 on both wave sizes, which collapses its two wave-specific check lines
into one and lets its Makefile stop selecting a prefix; and
workgroup_size_option2's 128-thread kernel goes from 129 to 192. The thirdcommit moves all seven to
smoke-dev, sincesmokeandsmoke-limboare notversioned and have to keep passing against the compilers already in the field,
which still produce the old numbers. That follows no-loop-4 and no-loop-7 in
aa1aa88; these can come back in twelve weeks.
Testing
gfx90a (MI210), against a build of ROCm/llvm-project#4196 at
c284a1582d8f:smoke-devafter the movesmokewas 321/321 andsmoke-limbo70/70The wave32 values are inferred rather than measured, as the machine has no
gfx10/11/12 part, so the gfx110X and gfx1151 legs of ROCm/llvm-project#4196 are
what will confirm them.
Made with Cursor