Repository navigation
fix(stinkytofu): add SAtomicCmpswapX2 to logical IR and rocisa adaptor shim - #11533
Conversation
…r shim Close AIHPBLAS-4621
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #11533 +/- ##
===========================================
- Coverage 69.34% 69.33% -0.01%
===========================================
Files 2817 2817
Lines 466948 467067 +119
Branches 68828 68858 +30
===========================================
+ Hits 323765 323814 +49
- Misses 119629 119677 +48
- Partials 23554 23576 +22
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
|
|
||
|
|
||
| # logicalIR: SAtomicCmpswapX2 | ||
| class SAtomicCmpswapX2(Instruction): |
There was a problem hiding this comment.
Thanks for picking this up — the SAtomicCmpswapX2 shim itself looks right (in particular getSrcParams() returning [base, soffset] matches native rocisa's AtomicReadWriteInstruction).
Heads-up though: this PR alone does not unblock the import. I applied it cleanly onto develop (6fdafda1d9a) and re-ran the failing import — it still raises, just pointing at a different class:
ImportError: cannot import name 'SAtomicUmaxX2' from 'rocisa_stinkytofu_adaptor.instruction'
The reason is that the SAtomicUmaxX2 instruction class sits on the same import statement as SAtomicCmpswapX2:
# Tensile/Components/SdmaRingEmitter.py:6-14
from rocisa.instruction import (
...
SStoreB32, SStoreB64, SStoreB128, SAtomicCmpswapX2, SAtomicUmaxX2,
)Python resolves that statement atomically, so one unresolvable class aborts the whole import regardless of the other 23 being present.
There is also a third instruction class missing on an adjacent path — SBfmB64, imported at module level by Tensile/Components/GlobalWriteBatch.py:32. Coverage after this PR:
| Instruction class | imported by | present after this PR |
|---|---|---|
SAtomicCmpswapX2 |
SdmaRingEmitter.py:13 |
✅ |
SAtomicUmaxX2 |
SdmaRingEmitter.py:13 |
❌ |
SBfmB64 |
GlobalWriteBatch.py:32 |
❌ |
All three came in together with #10925 (59f8fd37e86), which landed one day after #11485 (13d7206ad42) made gfx1250 auto-select the adaptor backend — so the adaptor never grew the corresponding shim classes. Worth noting the blast radius: Signature.py:62 imports SdmaRingEmitter unconditionally, so this breaks every gfx1250 kernel, not just the fused-A2A path.
Would you be open to folding the remaining two into this PR? I have both shims written and verified locally — happy to hand them over or push the delta to your branch, whichever is easier for you.
fix(stinkytofu): add SAtomicCmpswapX2 to logical IR and rocisa adaptor shim (#11533) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit JIRA ID : AIHPBLAS-4621 ## Motivation `from rocisa.instruction import SAtomicCmpswapX2` in `SdmaRingEmitter.py` raised `ImportError: cannot import name 'SAtomicCmpswapX2' from 'rocisa_stinkytofu_adaptor.instruction'`. Real rocisa's C++ backend already implements `s_atomic_cmpswap_x2` (`rocisa/rocisa/include/instruction/mem.hpp`, exercised by `rocisa/test/test_atomics.py`), but the gfx1250 StinkyTofu path was missing it in two places: stinkytofu's own logical IR (`LogicalInstructionDefs.inc`) never defined `SAtomicCmpswapX2`, and the Python adaptor shim (`rocisa_stinkytofu_adaptor/instruction.py`) that stands in for rocisa on gfx1250 never mirrored it, unlike its siblings `SAtomicInc`/`SAtomicDec`. ## Technical Details - `shared/stinkytofu/src/ir/logical/LogicalInstructionDefs.inc`: added an `SAtomicCmpswapX2` entry (numSrcs=2, hasDest=true, category "Memory Scalar", flag `IF_SMemAtomic`), mirroring the existing `SAtomicInc`/`SAtomicDec` entries. Rebuilding runs tablegen, which regenerates the C++ IR class and the nanobind Python bindings from this single definition. - `shared/stinkytofu/tests/unit/logical/LogicalToAsmMultiArchTest.cpp`: added the matching factory case and coverage-list entry so the new opcode is exercised by the exhaustive logical-IR test. - `projects/hipblaslt/tensilelite/rocisa_stinkytofu_adaptor/rocisa_stinkytofu_adaptor/instruction.py`: added an `SAtomicCmpswapX2` shim class mirroring `SAtomicInc`/`SAtomicDec` (`s_atomic_cmpswap_x2 dst, base, soffset`, with `to_stinky_logical()` forwarding to the new `stinkytofu.SAtomicCmpswapX2`). ## Test Plan - Rebuilt stinkytofu (`invoke build`) to regenerate tablegen output; confirmed `stinkytofu.SAtomicCmpswapX2` is now importable from the built Python module. - `./tests/unit_tests --gtest_filter="*LogicalToAsm*"` — all pass, 369/369 factory cases covered. - `pytest rocisa_stinkytofu_adaptor/tests/test_instruction.py` — 268 passed / 785 subtests passed (no regressions). - Manually instantiated the new adaptor shim class end-to-end: string rendering matches real rocisa's `test_atomics.py` expected output (`s_atomic_cmpswap_x2 s[4:7], s[8:9], s12 offset:64 glc`), `to_stinky_logical()` produces a valid `stinkytofu.LogicalInstruction`, and `deepcopy` round-trips. ## Test Result All of the above passed with no failures or regressions. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. ## Risk level Low — additive-only change (new instruction definition + new shim class), no existing behavior modified; verified against the existing test suites above.
…tor (#11592) JIRA ID : AIHPBLAS-4621 ## Motivation PR [#11178](#11178) added scalar atomic instructions (`SAtomicUmaxX2`, `SAtomicCmpswapX2`) and `SBfmB64` to native rocisa. A follow-up PR [#11533](#11533) added `SAtomicCmpswapX2` to the logical IR and adaptor, but missed `SAtomicUmaxX2` and `SBfmB64`. When the stinkytofu backend is active on gfx1250, importing these instructions from `rocisa.instruction` fails with `ImportError`. <!-- Explain the purpose of this PR and the goals it aims to achieve. --> ## Technical Details - **`SAtomicUmaxX2`**: Added logical IR definition in `LogicalInstructionDefs.inc` (numSrcs=2, hasDest=true, IF_SMemAtomic), full adaptor shim class with `to_stinky_logical()` in `instruction.py`, and test case in `LogicalToAsmMultiArchTest.cpp`. - **`SBfmB64`**: Added logical IR definition in `LogicalInstructionDefs.inc` (numSrcs=2, hasDest=true, IF_SALU), adaptor shim via `_make_scalar_alu_class` factory in `instruction.py`, and test case in `LogicalToAsmMultiArchTest.cpp`. <!-- Explain the changes along with any relevant GitHub links. --> ## Test Plan gfx1250 YAML test suites (gemm, sparse, streamk, gradient, streamk_ext). <!-- Explain any relevant testing done to verify this PR. --> ## Test Result All gfx1250 YAML tests pass. <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [V] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
fix(stinkytofu): add SAtomicUmaxX2 and SBfmB64 to logical IR and adaptor (#11592) JIRA ID : AIHPBLAS-4621 ## Motivation PR [#11178](ROCm/rocm-libraries#11178) added scalar atomic instructions (`SAtomicUmaxX2`, `SAtomicCmpswapX2`) and `SBfmB64` to native rocisa. A follow-up PR [#11533](ROCm/rocm-libraries#11533) added `SAtomicCmpswapX2` to the logical IR and adaptor, but missed `SAtomicUmaxX2` and `SBfmB64`. When the stinkytofu backend is active on gfx1250, importing these instructions from `rocisa.instruction` fails with `ImportError`. <!-- Explain the purpose of this PR and the goals it aims to achieve. --> ## Technical Details - **`SAtomicUmaxX2`**: Added logical IR definition in `LogicalInstructionDefs.inc` (numSrcs=2, hasDest=true, IF_SMemAtomic), full adaptor shim class with `to_stinky_logical()` in `instruction.py`, and test case in `LogicalToAsmMultiArchTest.cpp`. - **`SBfmB64`**: Added logical IR definition in `LogicalInstructionDefs.inc` (numSrcs=2, hasDest=true, IF_SALU), adaptor shim via `_make_scalar_alu_class` factory in `instruction.py`, and test case in `LogicalToAsmMultiArchTest.cpp`. <!-- Explain the changes along with any relevant GitHub links. --> ## Test Plan gfx1250 YAML test suites (gemm, sparse, streamk, gradient, streamk_ext). <!-- Explain any relevant testing done to verify this PR. --> ## Test Result All gfx1250 YAML tests pass. <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [V] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
TheRock Submodule Bump ActivityNewest first
|
JIRA ID : AIHPBLAS-4621
Motivation
from rocisa.instruction import SAtomicCmpswapX2inSdmaRingEmitter.pyraisedImportError: cannot import name 'SAtomicCmpswapX2' from 'rocisa_stinkytofu_adaptor.instruction'.Real rocisa's C++ backend already implements
s_atomic_cmpswap_x2(
rocisa/rocisa/include/instruction/mem.hpp, exercised byrocisa/test/test_atomics.py),but the gfx1250 StinkyTofu path was missing it in two places: stinkytofu's own
logical IR (
LogicalInstructionDefs.inc) never definedSAtomicCmpswapX2, and thePython adaptor shim (
rocisa_stinkytofu_adaptor/instruction.py) that stands in forrocisa on gfx1250 never mirrored it, unlike its siblings
SAtomicInc/SAtomicDec.Technical Details
shared/stinkytofu/src/ir/logical/LogicalInstructionDefs.inc: added anSAtomicCmpswapX2entry (numSrcs=2, hasDest=true, category "Memory Scalar",flag
IF_SMemAtomic), mirroring the existingSAtomicInc/SAtomicDecentries.Rebuilding runs tablegen, which regenerates the C++ IR class and the nanobind
Python bindings from this single definition.
shared/stinkytofu/tests/unit/logical/LogicalToAsmMultiArchTest.cpp: added thematching factory case and coverage-list entry so the new opcode is exercised by
the exhaustive logical-IR test.
projects/hipblaslt/tensilelite/rocisa_stinkytofu_adaptor/rocisa_stinkytofu_adaptor/instruction.py:added an
SAtomicCmpswapX2shim class mirroringSAtomicInc/SAtomicDec(
s_atomic_cmpswap_x2 dst, base, soffset, withto_stinky_logical()forwardingto the new
stinkytofu.SAtomicCmpswapX2).Test Plan
invoke build) to regenerate tablegen output; confirmedstinkytofu.SAtomicCmpswapX2is now importable from the built Python module../tests/unit_tests --gtest_filter="*LogicalToAsm*"— all pass, 369/369 factorycases covered.
pytest rocisa_stinkytofu_adaptor/tests/test_instruction.py— 268 passed / 785subtests passed (no regressions).
matches real rocisa's
test_atomics.pyexpected output(
s_atomic_cmpswap_x2 s[4:7], s[8:9], s12 offset:64 glc),to_stinky_logical()produces a valid
stinkytofu.LogicalInstruction, anddeepcopyround-trips.Test Result
All of the above passed with no failures or regressions.
Submission Checklist
Risk level
Low — additive-only change (new instruction definition + new shim class), no
existing behavior modified; verified against the existing test suites above.