From 7a67b2b9670328535d227f73051e710d59bd3e42 Mon Sep 17 00:00:00 2001
From: "Joao P. L. de Carvalho" <joao.carvalho@amd.com>
Date: Sun, 26 Apr 2026 20:42:41 +0000
Subject: [PATCH 1/8] Preserve MUBUF scope modifiers through StinkyTofu
 lowering

Carry rocisa MUBUF scope metadata through StinkyTofu conversion, IR
serialization, and assembly emission so scope-qualified buffer
operations survive lowering.
---
 ...{test_mubuf_off_vaddr.py => test_mubuf.py} | 59 ++++++++++++++++++-
 .../stinkytofu/ir/asm/StinkyModifiers.hpp     |  9 ++-
 .../conversion/rocisa/ToStinkyTofuUtils.cpp   |  6 +-
 .../serialization/asm/ModifierSerializer.cpp  |  7 ++-
 .../serialization/asm/StinkyAsmEmitter.cpp    |  3 +
 .../tests/filecheck/mubuf_off_vaddr.stir      | 10 ++++
 .../tests/unit/asm/StinkyAsmEmitterTest.cpp   | 30 ++++++++++
 7 files changed, 117 insertions(+), 7 deletions(-)
 rename projects/hipblaslt/tensilelite/rocisa/test/{test_mubuf_off_vaddr.py => test_mubuf.py} (65%)

diff --git a/projects/hipblaslt/tensilelite/rocisa/test/test_mubuf_off_vaddr.py b/projects/hipblaslt/tensilelite/rocisa/test/test_mubuf.py
similarity index 65%
rename from projects/hipblaslt/tensilelite/rocisa/test/test_mubuf_off_vaddr.py
rename to projects/hipblaslt/tensilelite/rocisa/test/test_mubuf.py
index 4a0d4d94411..2b62e277238 100644
--- a/projects/hipblaslt/tensilelite/rocisa/test/test_mubuf_off_vaddr.py
+++ b/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,
+    )
diff --git a/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp b/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp
index b9d360bb21c..c835f8bc40c 100644
--- a/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp
+++ b/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> {
diff --git a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
index 338e8de5634..3877efc9789 100644
--- a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
+++ b/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,
diff --git a/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp b/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp
index 623ad6aa92c..73889af4565 100644
--- a/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp
+++ b/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;
 }
@@ -380,7 +383,9 @@ void deserializeVisit(StinkyInstruction* inst, const std::string& attrKey,
         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)));
+                           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)));
diff --git a/shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp b/shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp
index 43761f6bd22..16fcc193577 100644
--- a/shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp
+++ b/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";
     }
diff --git a/shared/stinkytofu/tests/filecheck/mubuf_off_vaddr.stir b/shared/stinkytofu/tests/filecheck/mubuf_off_vaddr.stir
index 8669ecb2ac8..84c4276be43 100644
--- a/shared/stinkytofu/tests/filecheck/mubuf_off_vaddr.stir
+++ b/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" } }
+}
diff --git a/shared/stinkytofu/tests/unit/asm/StinkyAsmEmitterTest.cpp b/shared/stinkytofu/tests/unit/asm/StinkyAsmEmitterTest.cpp
index 17182231430..2ed655f42ed 100644
--- a/shared/stinkytofu/tests/unit/asm/StinkyAsmEmitterTest.cpp
+++ b/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);
+}

From 089ba4670c4be16cc3e50664773d90e47aa99637 Mon Sep 17 00:00:00 2001
From: "Joao P. L. de Carvalho" <joao.carvalho@amd.com>
Date: Mon, 27 Apr 2026 10:24:42 -0500
Subject: [PATCH 2/8] Format fix

---
 .../src/serialization/asm/ModifierSerializer.cpp      | 11 +++++------
 1 file changed, 5 insertions(+), 6 deletions(-)

diff --git a/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp b/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp
index 73889af4565..ec1c4530444 100644
--- a/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp
+++ b/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp
@@ -380,12 +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),
-                           false, false, false, false, getStr(fields, "scope", "") != "",
-                           getStr(fields, "scope", "")));
+        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)));

From e79e255f1edbc5a6383ca0ac6e01b77316e63a1e Mon Sep 17 00:00:00 2001
From: "Joao P. L. de Carvalho" <joao.carvalho@amd.com>
Date: Mon, 27 Apr 2026 18:41:38 -0500
Subject: [PATCH 3/8] Use enun rather than std::string for MUBUF modifiers

---
 .../stinkytofu/ir/asm/StinkyModifiers.hpp     | 36 +++++++++++++++++--
 .../conversion/rocisa/ToStinkyTofuUtils.cpp   | 18 ++++++++--
 .../serialization/asm/ModifierSerializer.cpp  | 15 ++++----
 .../serialization/asm/StinkyAsmEmitter.cpp    |  4 +--
 .../tests/unit/asm/StinkyAsmEmitterTest.cpp   |  2 +-
 5 files changed, 61 insertions(+), 14 deletions(-)

diff --git a/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp b/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp
index c835f8bc40c..d7db3688491 100644
--- a/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp
+++ b/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp
@@ -26,6 +26,7 @@
 #include <cstring>
 #include <memory>
 #include <string>
+#include <string_view>
 #include <vector>
 
 namespace stinkytofu {
@@ -34,6 +35,37 @@ enum class HighBitSel : int { NONE = -1, LOW = 0, HIGH = 1 };
 
 enum class MatrixFmt : uint8_t { FP4 = 0, FP6 = 1, FP8 = 2 };
 
+enum class MUBUFScope : uint8_t {
+    SCOPE_NONE = 0,
+    SCOPE_CU = 1,
+    SCOPE_SE = 2,
+    SCOPE_DEV = 3,
+    SCOPE_SYS = 4
+};
+
+inline std::string_view toString(MUBUFScope scope) {
+    switch (scope) {
+        case MUBUFScope::SCOPE_CU:
+            return "SCOPE_CU";
+        case MUBUFScope::SCOPE_SE:
+            return "SCOPE_SE";
+        case MUBUFScope::SCOPE_DEV:
+            return "SCOPE_DEV";
+        case MUBUFScope::SCOPE_SYS:
+            return "SCOPE_SYS";
+        default:
+            return "";
+    }
+}
+
+inline MUBUFScope parseMUBUFScope(std::string_view scope) {
+    if (scope == "SCOPE_CU") return MUBUFScope::SCOPE_CU;
+    if (scope == "SCOPE_SE") return MUBUFScope::SCOPE_SE;
+    if (scope == "SCOPE_DEV") return MUBUFScope::SCOPE_DEV;
+    if (scope == "SCOPE_SYS") return MUBUFScope::SCOPE_SYS;
+    return MUBUFScope::SCOPE_NONE;
+}
+
 struct Modifier {
     enum class Type : uint8_t {
         DS,
@@ -152,7 +184,7 @@ struct MUBUFModifiers : public TypedModifier<MUBUFModifiers> {
                    bool nt = false, bool lds = false, bool isStore = false,
                    bool hasMUBUFConst = false, bool hasGLCModifier = false,
                    bool hasSC0Modifier = false, bool hasSCOPEModifier = false,
-                   const std::string& scope = "")
+                   MUBUFScope scope = MUBUFScope::SCOPE_NONE)
         : TypedModifier<MUBUFModifiers>(),
           offset12(offset12),
           scope(scope),
@@ -168,7 +200,7 @@ struct MUBUFModifiers : public TypedModifier<MUBUFModifiers> {
           hasSCOPEModifier(hasSCOPEModifier) {}
 
     int offset12;
-    std::string scope;
+    MUBUFScope scope;
     uint32_t offen : 1;
     uint32_t glc : 1;
     uint32_t slc : 1;
diff --git a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
index 3877efc9789..1763da75e69 100644
--- a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
+++ b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
@@ -80,14 +80,28 @@ stinkytofu::FLATModifiers convertFLATModifiers(const rocisa::FLATModifiers& rocM
                                      rocMod.isStore, hasGLCModifier, hasSC0Modifier);
 }
 
+stinkytofu::MUBUFScope convertMUBUFScope(rocisa::CacheScope scope) {
+    switch (scope) {
+        case rocisa::CacheScope::SCOPE_CU:
+            return stinkytofu::MUBUFScope::SCOPE_CU;
+        case rocisa::CacheScope::SCOPE_SE:
+            return stinkytofu::MUBUFScope::SCOPE_SE;
+        case rocisa::CacheScope::SCOPE_DEV:
+            return stinkytofu::MUBUFScope::SCOPE_DEV;
+        case rocisa::CacheScope::SCOPE_SYS:
+            return stinkytofu::MUBUFScope::SCOPE_SYS;
+        default:
+            return stinkytofu::MUBUFScope::SCOPE_NONE;
+    }
+}
+
 stinkytofu::MUBUFModifiers convertMUBUFModifiers(const rocisa::MUBUFModifiers& rocMod,
                                                  const std::map<std::string, int>& asmCaps) {
     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);
+    stinkytofu::MUBUFScope scope = convertMUBUFScope(rocMod.scope);
     return stinkytofu::MUBUFModifiers(rocMod.offen, rocMod.offset12, rocMod.glc, rocMod.slc,
                                       rocMod.nt, rocMod.lds, rocMod.isStore, hasMUBUFConst,
                                       hasGLCModifier, hasSC0Modifier, hasSCOPEModifier, scope);
diff --git a/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp b/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp
index ec1c4530444..750d96ea6da 100644
--- a/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp
+++ b/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp
@@ -159,8 +159,8 @@ 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 << "\"";
+    if (mod.hasSCOPEModifier && mod.scope != MUBUFScope::SCOPE_NONE) {
+        os << ", scope = \"" << toString(mod.scope) << "\"";
     }
     os << " }";
     return true;
@@ -380,11 +380,12 @@ 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), false, false, false, false,
-            getStr(fields, "scope", "") != "", getStr(fields, "scope", "")));
+        MUBUFScope scope = parseMUBUFScope(getStr(fields, "scope", ""));
+        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, scope != MUBUFScope::SCOPE_NONE, scope));
     } else if (attrKey == "mod.smem") {
         inst->addModifier(SMEMModifiers(getBool(fields, "glc", false), getBool(fields, "nv", false),
                                         getInt(fields, "offset", 0)));
diff --git a/shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp b/shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp
index 16fcc193577..b8b49df1c87 100644
--- a/shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp
+++ b/shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp
@@ -203,8 +203,8 @@ 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.hasSCOPEModifier && mubufMod.scope != MUBUFScope::SCOPE_NONE) {
+        os << " scope:" << toString(mubufMod.scope);
     }
     if (mubufMod.nt) {
         os << " nt";
diff --git a/shared/stinkytofu/tests/unit/asm/StinkyAsmEmitterTest.cpp b/shared/stinkytofu/tests/unit/asm/StinkyAsmEmitterTest.cpp
index 2ed655f42ed..8623669e7c8 100644
--- a/shared/stinkytofu/tests/unit/asm/StinkyAsmEmitterTest.cpp
+++ b/shared/stinkytofu/tests/unit/asm/StinkyAsmEmitterTest.cpp
@@ -717,7 +717,7 @@ TEST_F(AsmEmitterTest, MUBUFScopeModifier) {
                             /*nt=*/false, /*lds=*/false, /*isStore=*/true,
                             /*hasMUBUFConst=*/false, /*hasGLCModifier=*/false,
                             /*hasSC0Modifier=*/false, /*hasSCOPEModifier=*/true,
-                            /*scope=*/"SCOPE_DEV");
+                            /*scope=*/MUBUFScope::SCOPE_DEV);
     inst->addModifier(mubufMod);
 
     AsmEmitterOptions options;

From d574f981d588af5cf4da6cf100afdab5b0dbc43c Mon Sep 17 00:00:00 2001
From: "Joao P. L. de Carvalho" <joao.carvalho@amd.com>
Date: Tue, 28 Apr 2026 20:03:28 +0000
Subject: [PATCH 4/8] Add offset flags if not set or no mubuf modifiers

---
 .../tensilelite/rocisa/rocisa/include/instruction/mem.hpp | 8 ++++++++
 1 file changed, 8 insertions(+)

diff --git a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/mem.hpp b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/mem.hpp
index 486cee166b0..f5d55deb506 100644
--- a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/mem.hpp
+++ b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/mem.hpp
@@ -352,6 +352,10 @@ namespace rocisa
             {
                 kStr += mubuf->toString();
             }
+            if(!mubuf || !mubuf->offen)
+            {
+                kStr += " offen offset:0";
+            }
             kStr = formatWithComment(kStr);
             setMsb(kStr, {vaddr}, dst);
             return kStr;
@@ -714,6 +718,10 @@ namespace rocisa
             {
                 kStr += mubuf->toString();
             }
+            if(!mubuf || !mubuf->offen)
+            {
+                kStr += " offen offset:0";
+            }
             kStr = formatWithComment(kStr);
             setMsb(kStr, {vaddr}, srcData);
             return kStr;

From 58c54c2041c5796b2b813b5b18c31de0c74195e4 Mon Sep 17 00:00:00 2001
From: "Henderson, Nathan" <Nathan.Henderson@amd.com>
Date: Wed, 29 Apr 2026 16:38:35 +0000
Subject: [PATCH 5/8] Fix MUBUF null soffset emission

Ensure real-vaddr MUBUF instructions keep an address modifier when literal
zero soffsets are lowered to null, while preserving the off-vaddr form.
---
 .../rocisa/rocisa/include/instruction/mem.hpp |  18 ++-
 .../tensilelite/rocisa/test/test_mubuf.py     | 135 +++++++++++++++++-
 .../stinkytofu/ir/asm/StinkyModifiers.hpp     |   6 +-
 .../conversion/rocisa/ToStinkyTofuUtils.cpp   |  44 +++++-
 4 files changed, 192 insertions(+), 11 deletions(-)

diff --git a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/mem.hpp b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/mem.hpp
index f5d55deb506..6c1e7a16391 100644
--- a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/mem.hpp
+++ b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/mem.hpp
@@ -275,6 +275,16 @@ namespace rocisa
         }
     };
 
+    /// Returns true when vaddr is the MUBUF "off" keyword.
+    inline bool isOffVAddr(const std::shared_ptr<Container>& vaddr)
+    {
+        if(auto* regCont = dynamic_cast<const RegisterContainer*>(vaddr.get()))
+        {
+            return regCont->isOff;
+        }
+        return false;
+    }
+
     struct MUBUFReadInstruction : public GlobalReadInstruction
     {
         std::shared_ptr<Container>    vaddr;
@@ -352,7 +362,7 @@ namespace rocisa
             {
                 kStr += mubuf->toString();
             }
-            if(!mubuf || !mubuf->offen)
+            if((!mubuf || !mubuf->offen) && !isOffVAddr(vaddr))
             {
                 kStr += " offen offset:0";
             }
@@ -718,7 +728,7 @@ namespace rocisa
             {
                 kStr += mubuf->toString();
             }
-            if(!mubuf || !mubuf->offen)
+            if((!mubuf || !mubuf->offen) && !isOffVAddr(vaddr))
             {
                 kStr += " offen offset:0";
             }
@@ -1624,6 +1634,10 @@ namespace rocisa
             std::string kStr = instStr + " " + getArgStr();
             if(mubuf)
                 kStr += mubuf->toString();
+            if((!mubuf || !mubuf->offen) && !isOffVAddr(vaddr))
+            {
+                kStr += " offen offset:0";
+            }
             kStr = formatWithComment(kStr);
             setMsb(kStr, {vaddr}, srcData);
             return kStr;
diff --git a/projects/hipblaslt/tensilelite/rocisa/test/test_mubuf.py b/projects/hipblaslt/tensilelite/rocisa/test/test_mubuf.py
index 2b62e277238..c5689246d0e 100644
--- a/projects/hipblaslt/tensilelite/rocisa/test/test_mubuf.py
+++ b/projects/hipblaslt/tensilelite/rocisa/test/test_mubuf.py
@@ -29,7 +29,13 @@
 from rocisa.code import Module, SignatureBase
 from rocisa.container import MUBUFModifiers, sgpr, vgpr
 from rocisa.enum import CacheScope
-from rocisa.instruction import BufferLoadB32, BufferStoreB32
+from rocisa.instruction import (
+    BufferAtomicAddF32,
+    BufferLoadB32,
+    BufferLoadB64,
+    BufferLoadB128,
+    BufferStoreB32,
+)
 
 _ISA = (12, 5, 0)
 
@@ -60,6 +66,14 @@ def _mubuf_off_asm() -> str:
             soffset=sgpr(46),
         )
     )
+    mod.add(
+        BufferStoreB32(
+            src=vgpr(13),
+            vaddr=vgpr("off", isOff=True),
+            saddr=sgpr(64, 4),
+            soffset=0,
+        )
+    )
     mod.setParent()  # resolves symbolic register names before conversion
 
     sig = SignatureBase(
@@ -80,8 +94,8 @@ 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)
+    assert re.search(r"buffer_store_b32 v13, off, s\[64:67\], null", _mubuf_off_asm)
 
 
 @pytest.fixture(scope="module")
@@ -134,3 +148,120 @@ def test_mubuf_scope_modifiers_stinkytofu(_mubuf_scope_asm):
         r"buffer_load_b32 v13, v33, s\[64:67\], s47 offen offset:0 scope:SCOPE_DEV",
         _mubuf_scope_asm,
     )
+
+
+def test_rocisa_load_null_soffset_adds_offen():
+    inst = BufferLoadB32(
+        dst=vgpr(13),
+        vaddr=vgpr(33),
+        saddr=sgpr(64, 4),
+        soffset=0,
+    )
+
+    assert str(inst).strip() == "buffer_load_b32 v13, v33, s[64:67], null offen offset:0"
+
+
+def test_rocisa_store_null_soffset_adds_offen():
+    inst = BufferStoreB32(
+        src=vgpr(12),
+        vaddr=vgpr(32),
+        saddr=sgpr(60, 4),
+        soffset=0,
+    )
+
+    assert str(inst).strip() == "buffer_store_b32 v12, v32, s[60:63], null offen offset:0"
+
+
+def test_rocisa_atomic_null_soffset_adds_offen():
+    inst = BufferAtomicAddF32(
+        src=vgpr(12),
+        vaddr=vgpr(32),
+        saddr=sgpr(60, 4),
+        soffset=0,
+    )
+
+    assert str(inst).strip() == "buffer_atomic_add_f32 v12, v32, s[60:63], null offen offset:0"
+
+
+def test_rocisa_off_vaddr_null_soffset():
+    inst = BufferStoreB32(
+        src=vgpr(12),
+        vaddr=vgpr("off", isOff=True),
+        saddr=sgpr(60, 4),
+        soffset=0,
+    )
+
+    assert str(inst).strip() == "buffer_store_b32 v12, off, s[60:63], null"
+
+
+@pytest.fixture(scope="module")
+def _mubuf_zero_soffset_asm() -> str:
+    mod = Module("mubuf_zero_soffset")
+    mod.add(
+        BufferLoadB128(
+            dst=vgpr(0, 4),
+            vaddr=vgpr(32),
+            saddr=sgpr(60, 4),
+            soffset=0,
+            mubuf=MUBUFModifiers(offen=True),
+            comment="G -> Reg 0_0_0_0",
+        )
+    )
+    mod.add(
+        BufferLoadB64(
+            dst=vgpr(4, 2),
+            vaddr=vgpr(33),
+            saddr=sgpr(60, 4),
+            soffset=0,
+            comment="G -> Reg 0_0_1_0",
+        )
+    )
+    mod.add(
+        BufferLoadB32(
+            dst=vgpr(6),
+            vaddr=vgpr(34),
+            saddr=sgpr(60, 4),
+            soffset=0,
+            mubuf=MUBUFModifiers(offen=False),
+            comment="G -> Reg 0_0_2_0",
+        )
+    )
+    mod.setParent()
+
+    sig = SignatureBase(
+        kernelName="mubuf_zero_soffset",
+        kernArgsVersion=1,
+        codeObjectVersion="4",
+        groupSegmentSize=0,
+        sgprWorkGroup=(1, 1, 0),
+        vgprWorkItem=0,
+        flatWorkGroupSize=64,
+        preloadKernArgs=False,
+    )
+
+    st = rocisa.toStinkyTofuModule(
+        mod, _ISA, "mubuf_zero_soffset", signature=sig, options={"OptLevel": 0}
+    )
+    st.runOptimizationPipeline()
+    return st.emitAssembly()
+
+
+def test_stinky_null_soffset_explicit_offen(_mubuf_zero_soffset_asm):
+    assert re.search(
+        r"buffer_load_b128 v\[0:3\], v32, s\[60:63\], null offen offset:0",
+        _mubuf_zero_soffset_asm,
+    )
+
+
+def test_stinky_null_soffset_default_mubuf(_mubuf_zero_soffset_asm):
+    assert re.search(
+        r"buffer_load_b64 v\[4:5\], v33, s\[60:63\], null offen offset:0",
+        _mubuf_zero_soffset_asm,
+    )
+
+
+def test_stinky_null_soffset_offen_false(_mubuf_zero_soffset_asm):
+    assert re.search(
+        r"buffer_load_b32 v6, v34, s\[60:63\], null offen offset:0",
+        _mubuf_zero_soffset_asm,
+    )
diff --git a/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp b/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp
index d7db3688491..2c7ff4b7272 100644
--- a/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp
+++ b/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp
@@ -187,7 +187,6 @@ struct MUBUFModifiers : public TypedModifier<MUBUFModifiers> {
                    MUBUFScope scope = MUBUFScope::SCOPE_NONE)
         : TypedModifier<MUBUFModifiers>(),
           offset12(offset12),
-          scope(scope),
           offen(offen),
           glc(glc),
           slc(slc),
@@ -197,10 +196,10 @@ struct MUBUFModifiers : public TypedModifier<MUBUFModifiers> {
           hasMUBUFConst(hasMUBUFConst),
           hasGLCModifier(hasGLCModifier),
           hasSC0Modifier(hasSC0Modifier),
-          hasSCOPEModifier(hasSCOPEModifier) {}
+          hasSCOPEModifier(hasSCOPEModifier),
+          scope(scope) {}
 
     int offset12;
-    MUBUFScope scope;
     uint32_t offen : 1;
     uint32_t glc : 1;
     uint32_t slc : 1;
@@ -211,6 +210,7 @@ struct MUBUFModifiers : public TypedModifier<MUBUFModifiers> {
     uint32_t hasGLCModifier : 1;
     uint32_t hasSC0Modifier : 1;
     uint32_t hasSCOPEModifier : 1;
+    MUBUFScope scope;
 };
 
 struct SMEMModifiers : public TypedModifier<SMEMModifiers> {
diff --git a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
index 1763da75e69..ad683695f1b 100644
--- a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
+++ b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
@@ -107,6 +107,38 @@ stinkytofu::MUBUFModifiers convertMUBUFModifiers(const rocisa::MUBUFModifiers& r
                                       hasGLCModifier, hasSC0Modifier, hasSCOPEModifier, scope);
 }
 
+/// Returns true when vaddr is the MUBUF "off" keyword.
+bool isOffVaddrContainer(const rocisa::Container* vaddr) {
+    if (auto* regCont = dynamic_cast<const rocisa::RegisterContainer*>(vaddr)) {
+        return regCont->isOff;
+    }
+    return false;
+}
+
+/// Build modifiers matching rocisa's MUBUF address form: real vaddr operands
+/// use `offen`, while `off` keeps the no-vaddr form.
+stinkytofu::MUBUFModifiers buildMUBUFModifiersForBufferOp(
+    const std::optional<rocisa::MUBUFModifiers>& rocMubuf, const rocisa::Container* vaddr,
+    const std::map<std::string, int>& asmCaps) {
+    bool hasMUBUFConst = asmCaps.count("HasMUBUFConst") && asmCaps.at("HasMUBUFConst");
+    bool hasGLCModifier = asmCaps.count("HasGLCModifier") && asmCaps.at("HasGLCModifier");
+    bool hasSC0Modifier = asmCaps.count("HasSC0Modifier") && asmCaps.at("HasSC0Modifier");
+
+    stinkytofu::MUBUFModifiers mod = rocMubuf.has_value()
+                                         ? convertMUBUFModifiers(rocMubuf.value(), asmCaps)
+                                         : stinkytofu::MUBUFModifiers(
+                                               /*offen=*/false, /*offset12=*/0, /*glc=*/false,
+                                               /*slc=*/false, /*nt=*/false, /*lds=*/false,
+                                               /*isStore=*/false, hasMUBUFConst, hasGLCModifier,
+                                               hasSC0Modifier, /*hasSCOPEModifier=*/false,
+                                               stinkytofu::MUBUFScope::SCOPE_NONE);
+
+    if (!mod.offen && !isOffVaddrContainer(vaddr)) {
+        mod.offen = 1;
+    }
+    return mod;
+}
+
 stinkytofu::SMEMModifiers convertSMEMModifiers(const rocisa::SMEMModifiers& rocMod,
                                                const std::map<std::string, int>& asmCaps) {
     bool hasSCOPEModifier = asmCaps.count("HasSCOPEModifier") && asmCaps.at("HasSCOPEModifier");
@@ -571,10 +603,14 @@ void addModifiersToInstruction(StinkyInstruction* stinkyInst, const rocisa::Inst
             [&](const auto& mod) { return convertFLATModifiers(mod, asmCaps); })
         else TRY_ADD_MOD(FLATStoreInstruction, flat, stinkytofu::FLATModifiers,
             [&](const auto& mod) { return convertFLATModifiers(mod, asmCaps); })
-        else TRY_ADD_MOD(MUBUFReadInstruction, mubuf, stinkytofu::MUBUFModifiers,
-            [&](const auto& mod) { return convertMUBUFModifiers(mod, asmCaps); })
-        else TRY_ADD_MOD(MUBUFStoreInstruction, mubuf, stinkytofu::MUBUFModifiers,
-            [&](const auto& mod) { return convertMUBUFModifiers(mod, asmCaps); })
+        else if (auto typed = dynamic_cast<const MUBUFReadInstruction*>(inst)) {
+            stinkyInst->addModifier<stinkytofu::MUBUFModifiers>(
+                buildMUBUFModifiersForBufferOp(typed->mubuf, typed->vaddr.get(), asmCaps));
+        }
+        else if (auto typed = dynamic_cast<const MUBUFStoreInstruction*>(inst)) {
+            stinkyInst->addModifier<stinkytofu::MUBUFModifiers>(
+                buildMUBUFModifiersForBufferOp(typed->mubuf, typed->vaddr.get(), asmCaps));
+        }
         else TRY_ADD_MOD(SMemLoadInstruction, smem, stinkytofu::SMEMModifiers,
             [&](const auto& mod) { return convertSMEMModifiers(mod, asmCaps); })
         else TRY_ADD_MOD(SMemStoreInstruction, smem, stinkytofu::SMEMModifiers,

From 676fad726dd61f0b6e3be8006c1579f6950c5394 Mon Sep 17 00:00:00 2001
From: "Henderson, Nathan" <Nathan.Henderson@amd.com>
Date: Wed, 29 Apr 2026 17:03:58 +0000
Subject: [PATCH 6/8] Fix format

---
 .../src/conversion/rocisa/ToStinkyTofuUtils.cpp   | 15 +++++++--------
 1 file changed, 7 insertions(+), 8 deletions(-)

diff --git a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
index ad683695f1b..e5d058c3a1c 100644
--- a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
+++ b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
@@ -124,14 +124,13 @@ stinkytofu::MUBUFModifiers buildMUBUFModifiersForBufferOp(
     bool hasGLCModifier = asmCaps.count("HasGLCModifier") && asmCaps.at("HasGLCModifier");
     bool hasSC0Modifier = asmCaps.count("HasSC0Modifier") && asmCaps.at("HasSC0Modifier");
 
-    stinkytofu::MUBUFModifiers mod = rocMubuf.has_value()
-                                         ? convertMUBUFModifiers(rocMubuf.value(), asmCaps)
-                                         : stinkytofu::MUBUFModifiers(
-                                               /*offen=*/false, /*offset12=*/0, /*glc=*/false,
-                                               /*slc=*/false, /*nt=*/false, /*lds=*/false,
-                                               /*isStore=*/false, hasMUBUFConst, hasGLCModifier,
-                                               hasSC0Modifier, /*hasSCOPEModifier=*/false,
-                                               stinkytofu::MUBUFScope::SCOPE_NONE);
+    stinkytofu::MUBUFModifiers mod =
+        rocMubuf.has_value() ? convertMUBUFModifiers(rocMubuf.value(), asmCaps)
+                             : stinkytofu::MUBUFModifiers(
+                                   /*offen=*/false, /*offset12=*/0, /*glc=*/false,
+                                   /*slc=*/false, /*nt=*/false, /*lds=*/false,
+                                   /*isStore=*/false, hasMUBUFConst, hasGLCModifier, hasSC0Modifier,
+                                   /*hasSCOPEModifier=*/false, stinkytofu::MUBUFScope::SCOPE_NONE);
 
     if (!mod.offen && !isOffVaddrContainer(vaddr)) {
         mod.offen = 1;

From aa9b63c9a2a00d01626a4bbd8bd273831178db84 Mon Sep 17 00:00:00 2001
From: "Henderson, Nathan" <Nathan.Henderson@amd.com>
Date: Thu, 30 Apr 2026 15:08:02 +0000
Subject: [PATCH 7/8] Remove redundant MUBUF scope flag

---
 .../stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp | 5 +----
 .../stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp   | 5 ++---
 .../stinkytofu/src/serialization/asm/ModifierSerializer.cpp  | 4 ++--
 shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp | 2 +-
 shared/stinkytofu/tests/unit/asm/StinkyAsmEmitterTest.cpp    | 3 +--
 5 files changed, 7 insertions(+), 12 deletions(-)

diff --git a/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp b/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp
index 2c7ff4b7272..d6164a25987 100644
--- a/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp
+++ b/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp
@@ -183,8 +183,7 @@ 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 hasSCOPEModifier = false,
-                   MUBUFScope scope = MUBUFScope::SCOPE_NONE)
+                   bool hasSC0Modifier = false, MUBUFScope scope = MUBUFScope::SCOPE_NONE)
         : TypedModifier<MUBUFModifiers>(),
           offset12(offset12),
           offen(offen),
@@ -196,7 +195,6 @@ struct MUBUFModifiers : public TypedModifier<MUBUFModifiers> {
           hasMUBUFConst(hasMUBUFConst),
           hasGLCModifier(hasGLCModifier),
           hasSC0Modifier(hasSC0Modifier),
-          hasSCOPEModifier(hasSCOPEModifier),
           scope(scope) {}
 
     int offset12;
@@ -209,7 +207,6 @@ struct MUBUFModifiers : public TypedModifier<MUBUFModifiers> {
     uint32_t hasMUBUFConst : 1;
     uint32_t hasGLCModifier : 1;
     uint32_t hasSC0Modifier : 1;
-    uint32_t hasSCOPEModifier : 1;
     MUBUFScope scope;
 };
 
diff --git a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
index e5d058c3a1c..2da0ffb511b 100644
--- a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
+++ b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
@@ -100,11 +100,10 @@ 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");
     stinkytofu::MUBUFScope scope = convertMUBUFScope(rocMod.scope);
     return stinkytofu::MUBUFModifiers(rocMod.offen, rocMod.offset12, rocMod.glc, rocMod.slc,
                                       rocMod.nt, rocMod.lds, rocMod.isStore, hasMUBUFConst,
-                                      hasGLCModifier, hasSC0Modifier, hasSCOPEModifier, scope);
+                                      hasGLCModifier, hasSC0Modifier, scope);
 }
 
 /// Returns true when vaddr is the MUBUF "off" keyword.
@@ -130,7 +129,7 @@ stinkytofu::MUBUFModifiers buildMUBUFModifiersForBufferOp(
                                    /*offen=*/false, /*offset12=*/0, /*glc=*/false,
                                    /*slc=*/false, /*nt=*/false, /*lds=*/false,
                                    /*isStore=*/false, hasMUBUFConst, hasGLCModifier, hasSC0Modifier,
-                                   /*hasSCOPEModifier=*/false, stinkytofu::MUBUFScope::SCOPE_NONE);
+                                   stinkytofu::MUBUFScope::SCOPE_NONE);
 
     if (!mod.offen && !isOffVaddrContainer(vaddr)) {
         mod.offen = 1;
diff --git a/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp b/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp
index 750d96ea6da..39388f9fdda 100644
--- a/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp
+++ b/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp
@@ -159,7 +159,7 @@ 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 != MUBUFScope::SCOPE_NONE) {
+    if (mod.scope != MUBUFScope::SCOPE_NONE) {
         os << ", scope = \"" << toString(mod.scope) << "\"";
     }
     os << " }";
@@ -385,7 +385,7 @@ void deserializeVisit(StinkyInstruction* inst, const std::string& attrKey,
             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, scope != MUBUFScope::SCOPE_NONE, scope));
+                           false, false, false, scope));
     } else if (attrKey == "mod.smem") {
         inst->addModifier(SMEMModifiers(getBool(fields, "glc", false), getBool(fields, "nv", false),
                                         getInt(fields, "offset", 0)));
diff --git a/shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp b/shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp
index b8b49df1c87..c4884623a38 100644
--- a/shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp
+++ b/shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp
@@ -203,7 +203,7 @@ inline std::ostream& operator<<(std::ostream& os, const MUBUFModifiers& mubufMod
         else if (mubufMod.hasSC0Modifier)
             os << " sc1";
     }
-    if (mubufMod.hasSCOPEModifier && mubufMod.scope != MUBUFScope::SCOPE_NONE) {
+    if (mubufMod.scope != MUBUFScope::SCOPE_NONE) {
         os << " scope:" << toString(mubufMod.scope);
     }
     if (mubufMod.nt) {
diff --git a/shared/stinkytofu/tests/unit/asm/StinkyAsmEmitterTest.cpp b/shared/stinkytofu/tests/unit/asm/StinkyAsmEmitterTest.cpp
index 8623669e7c8..8bdcc69facc 100644
--- a/shared/stinkytofu/tests/unit/asm/StinkyAsmEmitterTest.cpp
+++ b/shared/stinkytofu/tests/unit/asm/StinkyAsmEmitterTest.cpp
@@ -716,8 +716,7 @@ TEST_F(AsmEmitterTest, MUBUFScopeModifier) {
     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=*/MUBUFScope::SCOPE_DEV);
+                            /*hasSC0Modifier=*/false, /*scope=*/MUBUFScope::SCOPE_DEV);
     inst->addModifier(mubufMod);
 
     AsmEmitterOptions options;

From 53bf38bda6bbb17124ea7796133c1073f6056e27 Mon Sep 17 00:00:00 2001
From: "Henderson, Nathan" <Nathan.Henderson@amd.com>
Date: Thu, 30 Apr 2026 15:23:26 +0000
Subject: [PATCH 8/8] Fix format

---
 .../src/conversion/rocisa/ToStinkyTofuUtils.cpp    | 14 +++++++-------
 1 file changed, 7 insertions(+), 7 deletions(-)

diff --git a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
index 2da0ffb511b..c7425f08781 100644
--- a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
+++ b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp
@@ -123,13 +123,13 @@ stinkytofu::MUBUFModifiers buildMUBUFModifiersForBufferOp(
     bool hasGLCModifier = asmCaps.count("HasGLCModifier") && asmCaps.at("HasGLCModifier");
     bool hasSC0Modifier = asmCaps.count("HasSC0Modifier") && asmCaps.at("HasSC0Modifier");
 
-    stinkytofu::MUBUFModifiers mod =
-        rocMubuf.has_value() ? convertMUBUFModifiers(rocMubuf.value(), asmCaps)
-                             : stinkytofu::MUBUFModifiers(
-                                   /*offen=*/false, /*offset12=*/0, /*glc=*/false,
-                                   /*slc=*/false, /*nt=*/false, /*lds=*/false,
-                                   /*isStore=*/false, hasMUBUFConst, hasGLCModifier, hasSC0Modifier,
-                                   stinkytofu::MUBUFScope::SCOPE_NONE);
+    stinkytofu::MUBUFModifiers mod = rocMubuf.has_value()
+                                         ? convertMUBUFModifiers(rocMubuf.value(), asmCaps)
+                                         : stinkytofu::MUBUFModifiers(
+                                               /*offen=*/false, /*offset12=*/0, /*glc=*/false,
+                                               /*slc=*/false, /*nt=*/false, /*lds=*/false,
+                                               /*isStore=*/false, hasMUBUFConst, hasGLCModifier,
+                                               hasSC0Modifier, stinkytofu::MUBUFScope::SCOPE_NONE);
 
     if (!mod.offen && !isOffVaddrContainer(vaddr)) {
         mod.offen = 1;