[stinkytofu] Preserve MUBUF scope modifiers through StinkyTofu lowering

projects/hipblaslt/tensilelite/rocisa/test/test_mubuf_off_vaddr.py → projects/hipblaslt/tensilelite/rocisa/test/test_mubuf.py

@@ -20,15 +20,16 @@
 # CTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
 ################################################################################
 
-"""Regression: MUBUF isOff=True vaddr must survive rocisa → StinkyTofu as 'off', not 'v[vgproff]'."""
+"""Regressions for MUBUF rocisa -> StinkyTofu lowering."""
 
 import re
 
 import pytest
 import rocisa
 from rocisa.code import Module, SignatureBase
-from rocisa.container import sgpr, vgpr
-from rocisa.instruction import BufferStoreB32
+from rocisa.container import MUBUFModifiers, sgpr, vgpr
+from rocisa.enum import CacheScope
+from rocisa.instruction import BufferLoadB32, BufferStoreB32
 
 _ISA = (12, 5, 0)
 
@@ -81,3 +82,55 @@ def _mubuf_off_asm() -> str:
 def test_mubuf_off_vaddr_stinkytofu(_mubuf_off_asm):
     # Assembler rejects 'v[vgproff]'; 'off' must appear as the literal vaddr operand.
     assert re.search(r"buffer_store_b32 v12, off, s\[60:63\], s46", _mubuf_off_asm)
+
+
+@pytest.fixture(scope="module")
+def _mubuf_scope_asm() -> str:
+    mod = Module("mubuf_scope_modifiers")
+    mod.add(
+        BufferStoreB32(
+            src=vgpr(12),
+            vaddr=vgpr(32),
+            saddr=sgpr(60, 4),
+            soffset=sgpr(46),
+            mubuf=MUBUFModifiers(offen=True, scope=CacheScope.SCOPE_DEV),
+        )
+    )
+    mod.add(
+        BufferLoadB32(
+            dst=vgpr(13),
+            vaddr=vgpr(33),
+            saddr=sgpr(64, 4),
+            soffset=sgpr(47),
+            mubuf=MUBUFModifiers(offen=True, scope=CacheScope.SCOPE_DEV),
+        )
+    )
+    mod.setParent()
+
+    sig = SignatureBase(
+        kernelName="mubuf_scope_modifiers",
+        kernArgsVersion=1,
+        codeObjectVersion="4",
+        groupSegmentSize=0,
+        sgprWorkGroup=(1, 1, 0),
+        vgprWorkItem=0,
+        flatWorkGroupSize=64,
+        preloadKernArgs=False,
+    )
+
+    st = rocisa.toStinkyTofuModule(
+        mod, _ISA, "mubuf_scope_modifiers", signature=sig, options={"OptLevel": 0}
+    )
+    st.runOptimizationPipeline()
+    return st.emitAssembly()
+
+
+def test_mubuf_scope_modifiers_stinkytofu(_mubuf_scope_asm):
+    assert re.search(
+        r"buffer_store_b32 v12, v32, s\[60:63\], s46 offen offset:0 scope:SCOPE_DEV",
+        _mubuf_scope_asm,
+    )
+    assert re.search(
+        r"buffer_load_b32 v13, v33, s\[64:67\], s47 offen offset:0 scope:SCOPE_DEV",
+        _mubuf_scope_asm,
+    )

shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp

@@ -151,9 +151,11 @@ struct MUBUFModifiers : public TypedModifier<MUBUFModifiers> {
     MUBUFModifiers(bool offen = false, int offset12 = 0, bool glc = false, bool slc = false,
                    bool nt = false, bool lds = false, bool isStore = false,
                    bool hasMUBUFConst = false, bool hasGLCModifier = false,
-                   bool hasSC0Modifier = false)
+                   bool hasSC0Modifier = false, bool hasSCOPEModifier = false,
+                   const std::string& scope = "")
         : TypedModifier<MUBUFModifiers>(),
           offset12(offset12),
+          scope(scope),
           offen(offen),
           glc(glc),
           slc(slc),
@@ -162,9 +164,11 @@ struct MUBUFModifiers : public TypedModifier<MUBUFModifiers> {
           isStore(isStore),
           hasMUBUFConst(hasMUBUFConst),
           hasGLCModifier(hasGLCModifier),
-          hasSC0Modifier(hasSC0Modifier) {}
+          hasSC0Modifier(hasSC0Modifier),
+          hasSCOPEModifier(hasSCOPEModifier) {}
 
     int offset12;
+    std::string scope;
     uint32_t offen : 1;
     uint32_t glc : 1;
     uint32_t slc : 1;
@@ -174,6 +178,7 @@ struct MUBUFModifiers : public TypedModifier<MUBUFModifiers> {
     uint32_t hasMUBUFConst : 1;
     uint32_t hasGLCModifier : 1;
     uint32_t hasSC0Modifier : 1;
+    uint32_t hasSCOPEModifier : 1;
 };
 
 struct SMEMModifiers : public TypedModifier<SMEMModifiers> {

shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp

@@ -31,6 +31,7 @@
 #include <cctype>
 #include <iostream>
 #include <optional>
+#include <string>
 #include <string_view>
 #include <vector>
 
@@ -84,9 +85,12 @@ stinkytofu::MUBUFModifiers convertMUBUFModifiers(const rocisa::MUBUFModifiers& r
     bool hasMUBUFConst = asmCaps.count("HasMUBUFConst") && asmCaps.at("HasMUBUFConst");
     bool hasGLCModifier = asmCaps.count("HasGLCModifier") && asmCaps.at("HasGLCModifier");
     bool hasSC0Modifier = asmCaps.count("HasSC0Modifier") && asmCaps.at("HasSC0Modifier");
+    bool hasSCOPEModifier = asmCaps.count("HasSCOPEModifier") && asmCaps.at("HasSCOPEModifier");
+    std::string scope =
+        rocMod.scope == rocisa::CacheScope::SCOPE_NONE ? "" : rocisa::toString(rocMod.scope);
     return stinkytofu::MUBUFModifiers(rocMod.offen, rocMod.offset12, rocMod.glc, rocMod.slc,
                                       rocMod.nt, rocMod.lds, rocMod.isStore, hasMUBUFConst,
-                                      hasGLCModifier, hasSC0Modifier);
+                                      hasGLCModifier, hasSC0Modifier, hasSCOPEModifier, scope);
 }
 
 stinkytofu::SMEMModifiers convertSMEMModifiers(const rocisa::SMEMModifiers& rocMod,

shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp

@@ -159,6 +159,9 @@ bool serializeVisit(const MUBUFModifiers& mod, std::ostream& os) {
     os << " offen = " << (mod.offen ? "true" : "false") << ", offset12 = " << mod.offset12
        << ", glc = " << (mod.glc ? "true" : "false") << ", slc = " << (mod.slc ? "true" : "false")
        << ", nt = " << (mod.nt ? "true" : "false") << ", lds = " << (mod.lds ? "true" : "false");
+    if (mod.hasSCOPEModifier && !mod.scope.empty()) {
+        os << ", scope = \"" << mod.scope << "\"";
+    }
     os << " }";
     return true;
 }
@@ -377,10 +380,11 @@ void deserializeVisit(StinkyInstruction* inst, const std::string& attrKey,
     } else if (attrKey == "mod.global") {
         inst->addModifier(GLOBALModifiers(getInt(fields, "offset", 0)));
     } else if (attrKey == "mod.mubuf") {
-        inst->addModifier(
-            MUBUFModifiers(getBool(fields, "offen", false), getInt(fields, "offset12", 0),
-                           getBool(fields, "glc", false), getBool(fields, "slc", false),
-                           getBool(fields, "nt", false), getBool(fields, "lds", false)));
+        inst->addModifier(MUBUFModifiers(
+            getBool(fields, "offen", false), getInt(fields, "offset12", 0),
+            getBool(fields, "glc", false), getBool(fields, "slc", false),
+            getBool(fields, "nt", false), getBool(fields, "lds", false), false, false, false, false,
+            getStr(fields, "scope", "") != "", getStr(fields, "scope", "")));
     } else if (attrKey == "mod.smem") {
         inst->addModifier(SMEMModifiers(getBool(fields, "glc", false), getBool(fields, "nv", false),
                                         getInt(fields, "offset", 0)));

shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp

@@ -203,6 +203,9 @@ inline std::ostream& operator<<(std::ostream& os, const MUBUFModifiers& mubufMod
         else if (mubufMod.hasSC0Modifier)
             os << " sc1";
     }
+    if (mubufMod.hasSCOPEModifier && !mubufMod.scope.empty()) {
+        os << " scope:" << mubufMod.scope;
+    }
     if (mubufMod.nt) {
         os << " nt";
     }

shared/stinkytofu/tests/filecheck/mubuf_off_vaddr.stir

@@ -25,3 +25,13 @@ st.func @buffer_load_b32_off_vaddr() {
 ^entry:
   v0 = "st.buffer_load_b32"(off, s[4:7], s3) { issueCycles = 1, latencyCycles = 100 }
 }
+
+#--- MUBUF scope modifier: IR round-trip must preserve scope for asm emission
+# CHECK-LABEL: @buffer_store_b32_scope
+# CHECK: mod.mubuf
+# CHECK-SAME: scope = "SCOPE_DEV"
+
+st.func @buffer_store_b32_scope() {
+^entry:
+  "st.buffer_store_b32"(v12, v32, s[60:63], s46) { issueCycles = 1, latencyCycles = 1, mod.mubuf = { offen = true, offset12 = 0, glc = false, slc = false, nt = false, lds = false, scope = "SCOPE_DEV" } }
+}

shared/stinkytofu/tests/unit/asm/StinkyAsmEmitterTest.cpp

@@ -703,3 +703,33 @@ TEST_F(AsmEmitterTest, MUBUFLoadB32_OffVAddr) {
     std::string expected = "buffer_load_b32 v0, off, s[4:7], s3\n";
     EXPECT_EQ(assembly, expected);
 }
+
+TEST_F(AsmEmitterTest, MUBUFScopeModifier) {
+    StinkyInstruction* inst = createInstruction("buffer_store_b32");
+    ASSERT_NE(inst, nullptr);
+
+    inst->addSrcReg(StinkyRegister("v", 12, 1));
+    inst->addSrcReg(StinkyRegister("v", 32, 1));
+    inst->addSrcReg(StinkyRegister("s", 60, 4));
+    inst->addSrcReg(StinkyRegister("s", 46, 1));
+
+    MUBUFModifiers mubufMod(/*offen=*/true, /*offset12=*/0, /*glc=*/false, /*slc=*/false,
+                            /*nt=*/false, /*lds=*/false, /*isStore=*/true,
+                            /*hasMUBUFConst=*/false, /*hasGLCModifier=*/false,
+                            /*hasSC0Modifier=*/false, /*hasSCOPEModifier=*/true,
+                            /*scope=*/"SCOPE_DEV");
+    inst->addModifier(mubufMod);
+
+    AsmEmitterOptions options;
+    options.emitComments = false;
+    options.emitCycleInfo = false;
+    options.indent = 0;
+    options.emitBlankLines = false;
+
+    StinkyAsmEmitter emitter(options);
+    std::string assembly = emitter.emit(*inst);
+
+    std::string expected =
+        "buffer_store_b32 v12, v32, s[60:63], s46 offen offset:0 scope:SCOPE_DEV\n";
+    EXPECT_EQ(assembly, expected);
+}