Skip to content

fix(stinkytofu): add SAtomicCmpswapX2 to logical IR and rocisa adaptor shim #11533

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
KKyang merged 1 commit into from
Sep 2, 2026
Merged

fix(stinkytofu): add SAtomicCmpswapX2 to logical IR and rocisa adaptor shim #11533

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
53 changes: 53 additions & 0 deletions .../hipblaslt/tensilelite/rocisa_stinkytofu_adaptor/rocisa_stinkytofu_adaptor/instruction.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4132,6 +4132,59 @@ def __deepcopy__(self, memo):
dup.smem = _deepcopy(self.smem, memo) if self.smem is not None else None
return dup


# logicalIR: SAtomicCmpswapX2
class SAtomicCmpswapX2(Instruction):

@jaopaulolc jaopaulolc Sep 1, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

"""``s_atomic_cmpswap_x2 dst, base, soffset`` shim."""

__slots__ = ("dst", "base", "soffset", "smem")

def __init__(self, dst=None, base=None, soffset=None, smem=None, comment="", **kw):
_ = kw
super().__init__(InstType.INST_B128, comment)
self.dst = dst
self.base = base
self.soffset = soffset
self.smem = smem
self.setInst("s_atomic_cmpswap_x2")

def getParams(self):
return [self.dst, self.base, self.soffset]

def getDstParams(self):
return [self.dst] if self.dst else []

def getSrcParams(self):
return [self.base, self.soffset]

def toString(self) -> str:
parts = [_input_to_str(self.dst), _input_to_str(self.base), _input_to_str(self.soffset)]
kstr = self.instStr + " " + ", ".join(parts)
if self.smem is not None and hasattr(self.smem, "toString"):
kstr += self.smem.toString()
return self.formatWithComment(kstr)

def to_stinky_logical(self) -> Any:
import stinkytofu as _st
return _st.SAtomicCmpswapX2(
_to_stinky_register(self.dst),
_to_stinky_register(self.base),
_to_stinky_register(self.soffset),
self.comment)

def __deepcopy__(self, memo):
if id(self) in memo:
return memo[id(self)]
dup = self.__class__.__new__(self.__class__)
memo[id(self)] = dup
Instruction.__init__(dup, self.instType, self.comment)
dup.instStr = self.instStr
dup.dst = _deepcopy(self.dst, memo) if self.dst is not None else None
dup.base = _deepcopy(self.base, memo) if self.base is not None else None
dup.soffset = self.soffset if isinstance(self.soffset, (int, float, str, bool)) else _deepcopy(self.soffset, memo)
dup.smem = _deepcopy(self.smem, memo) if self.smem is not None else None
return dup

# --- TensorLoadToLds: rocisa(group0, group1, group2, group3, comment) ---
def _make_tensor_load_class():
def __init__(self, group0: Any = None, group1: Any = None,
Expand Down
10 changes: 10 additions & 0 deletions shared/stinkytofu/src/ir/logical/LogicalInstructionDefs.inc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3254,6 +3254,16 @@
false,
false,
false,
"IF_SMemAtomic"},

{"SAtomicCmpswapX2",
"SAtomicCmpswapX2",
2,
true,
"Memory Scalar",
false,
false,
false,
"IF_SMemAtomic"},

{"SCSelectB64",
Expand Down
3 changes: 3 additions & 0 deletions shared/stinkytofu/tests/unit/logical/LogicalToAsmMultiArchTest.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -731,6 +731,8 @@ static LogicalInstruction* createTestInstruction(logical::Opcode opcode) {
return SAtomicInc(sgpr(0), sgpr(1), sgpr(2));
case logical::SAtomicDec:
return SAtomicDec(sgpr(0), sgpr(1));
case logical::SAtomicCmpswapX2:
return SAtomicCmpswapX2(sgpr(0, 4), sgpr(1), sgpr(2));
case logical::SCSelectB64:
return SCSelectB64(sgpr(0), sgpr(1), sgpr(2));
case logical::SCmpKEQU32:
Expand Down Expand Up @@ -1112,6 +1114,7 @@ TEST(LogicalToAsmComprehensive, AllInstructionsAllArchitectures) {
logical::SStoreB512,
logical::SAtomicInc,
logical::SAtomicDec,
logical::SAtomicCmpswapX2,
logical::SCSelectB64,
logical::SCmpKEQU32,
logical::SCmpKGeU32,
Expand Down
Loading