diff --git a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/mem.hpp b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/mem.hpp index 486cee166b0..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& vaddr) + { + if(auto* regCont = dynamic_cast(vaddr.get())) + { + return regCont->isOff; + } + return false; + } + struct MUBUFReadInstruction : public GlobalReadInstruction { std::shared_ptr vaddr; @@ -352,6 +362,10 @@ namespace rocisa { kStr += mubuf->toString(); } + if((!mubuf || !mubuf->offen) && !isOffVAddr(vaddr)) + { + kStr += " offen offset:0"; + } kStr = formatWithComment(kStr); setMsb(kStr, {vaddr}, dst); return kStr; @@ -714,6 +728,10 @@ namespace rocisa { kStr += mubuf->toString(); } + if((!mubuf || !mubuf->offen) && !isOffVAddr(vaddr)) + { + kStr += " offen offset:0"; + } kStr = formatWithComment(kStr); setMsb(kStr, {vaddr}, srcData); return kStr; @@ -1616,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 new file mode 100644 index 00000000000..c5689246d0e --- /dev/null +++ b/projects/hipblaslt/tensilelite/rocisa/test/test_mubuf.py @@ -0,0 +1,267 @@ +################################################################################ +# +# Copyright (C) 2025 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell cop- +# ies of the Software, and to permit persons to whom the Software is furnished +# to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in all +# copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IM- +# PLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS +# FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR +# COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER +# IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNE- +# CTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. +################################################################################ + +"""Regressions for MUBUF rocisa -> StinkyTofu lowering.""" + +import re + +import pytest +import rocisa +from rocisa.code import Module, SignatureBase +from rocisa.container import MUBUFModifiers, sgpr, vgpr +from rocisa.enum import CacheScope +from rocisa.instruction import ( + BufferAtomicAddF32, + BufferLoadB32, + BufferLoadB64, + BufferLoadB128, + BufferStoreB32, +) + +_ISA = (12, 5, 0) + +# Skip entire module when the target backend isn't compiled into the registry. +pytestmark = pytest.mark.skipif( + not rocisa.isSupportedByStinkyTofu(_ISA), + reason=f"gfx{''.join(str(v) for v in _ISA)} not registered in StinkyTofu BackendRegistry", +) + + +@pytest.fixture(scope="module", autouse=True) +def _isa_context(): + import os + + rocm_path = os.environ.get("ROCM_PATH", "/opt/rocm") + rocisa.rocIsa.getInstance().init(_ISA, rocm_path + "/bin/amdclang++", False) + rocisa.rocIsa.getInstance().setKernel(_ISA, 32) + + +@pytest.fixture(scope="module") +def _mubuf_off_asm() -> str: + mod = Module("mubuf_off_vaddr") + mod.add( + BufferStoreB32( + src=vgpr(12), + vaddr=vgpr("off", isOff=True), # isOff → MUBUF 'off' keyword, not a named VGPR + saddr=sgpr(60, 4), + 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( + kernelName="mubuf_off_vaddr", + kernArgsVersion=1, + codeObjectVersion="4", + groupSegmentSize=0, + sgprWorkGroup=(1, 1, 0), + vgprWorkItem=0, + flatWorkGroupSize=64, + preloadKernArgs=False, + ) + + stinky_module_options = {"OptLevel": 0} + st = rocisa.toStinkyTofuModule(mod, _ISA, "mubuf_off_vaddr", signature=sig, options=stinky_module_options) + st.runOptimizationPipeline() + return st.emitAssembly() + + +def test_mubuf_off_vaddr_stinkytofu(_mubuf_off_asm): + 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") +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, + ) + + +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/projects/hipblaslt/tensilelite/rocisa/test/test_mubuf_off_vaddr.py b/projects/hipblaslt/tensilelite/rocisa/test/test_mubuf_off_vaddr.py deleted file mode 100644 index 4a0d4d94411..00000000000 --- a/projects/hipblaslt/tensilelite/rocisa/test/test_mubuf_off_vaddr.py +++ /dev/null @@ -1,83 +0,0 @@ -################################################################################ -# -# Copyright (C) 2025 Advanced Micro Devices, Inc. All rights reserved. -# -# Permission is hereby granted, free of charge, to any person obtaining a copy -# of this software and associated documentation files (the "Software"), to deal -# in the Software without restriction, including without limitation the rights -# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell cop- -# ies of the Software, and to permit persons to whom the Software is furnished -# to do so, subject to the following conditions: -# -# The above copyright notice and this permission notice shall be included in all -# copies or substantial portions of the Software. -# -# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IM- -# PLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS -# FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR -# COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER -# IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNE- -# 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]'.""" - -import re - -import pytest -import rocisa -from rocisa.code import Module, SignatureBase -from rocisa.container import sgpr, vgpr -from rocisa.instruction import BufferStoreB32 - -_ISA = (12, 5, 0) - -# Skip entire module when the target backend isn't compiled into the registry. -pytestmark = pytest.mark.skipif( - not rocisa.isSupportedByStinkyTofu(_ISA), - reason=f"gfx{''.join(str(v) for v in _ISA)} not registered in StinkyTofu BackendRegistry", -) - - -@pytest.fixture(scope="module", autouse=True) -def _isa_context(): - import os - - rocm_path = os.environ.get("ROCM_PATH", "/opt/rocm") - rocisa.rocIsa.getInstance().init(_ISA, rocm_path + "/bin/amdclang++", False) - rocisa.rocIsa.getInstance().setKernel(_ISA, 32) - - -@pytest.fixture(scope="module") -def _mubuf_off_asm() -> str: - mod = Module("mubuf_off_vaddr") - mod.add( - BufferStoreB32( - src=vgpr(12), - vaddr=vgpr("off", isOff=True), # isOff → MUBUF 'off' keyword, not a named VGPR - saddr=sgpr(60, 4), - soffset=sgpr(46), - ) - ) - mod.setParent() # resolves symbolic register names before conversion - - sig = SignatureBase( - kernelName="mubuf_off_vaddr", - kernArgsVersion=1, - codeObjectVersion="4", - groupSegmentSize=0, - sgprWorkGroup=(1, 1, 0), - vgprWorkItem=0, - flatWorkGroupSize=64, - preloadKernArgs=False, - ) - - stinky_module_options = {"OptLevel": 0} - st = rocisa.toStinkyTofuModule(mod, _ISA, "mubuf_off_vaddr", signature=sig, options=stinky_module_options) - st.runOptimizationPipeline() - return st.emitAssembly() - - -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) diff --git a/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp b/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp index b9d360bb21c..d6164a25987 100644 --- a/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp +++ b/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp @@ -26,6 +26,7 @@ #include #include #include +#include #include 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, @@ -151,7 +183,7 @@ struct MUBUFModifiers : public TypedModifier { 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, MUBUFScope scope = MUBUFScope::SCOPE_NONE) : TypedModifier(), offset12(offset12), offen(offen), @@ -162,7 +194,8 @@ struct MUBUFModifiers : public TypedModifier { isStore(isStore), hasMUBUFConst(hasMUBUFConst), hasGLCModifier(hasGLCModifier), - hasSC0Modifier(hasSC0Modifier) {} + hasSC0Modifier(hasSC0Modifier), + scope(scope) {} int offset12; uint32_t offen : 1; @@ -174,6 +207,7 @@ struct MUBUFModifiers : public TypedModifier { uint32_t hasMUBUFConst : 1; uint32_t hasGLCModifier : 1; uint32_t hasSC0Modifier : 1; + MUBUFScope scope; }; struct SMEMModifiers : public TypedModifier { diff --git a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp index 338e8de5634..c7425f08781 100644 --- a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp +++ b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp @@ -31,6 +31,7 @@ #include #include #include +#include #include #include @@ -79,14 +80,61 @@ 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& 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::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); + hasGLCModifier, hasSC0Modifier, scope); +} + +/// Returns true when vaddr is the MUBUF "off" keyword. +bool isOffVaddrContainer(const rocisa::Container* vaddr) { + if (auto* regCont = dynamic_cast(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& rocMubuf, const rocisa::Container* vaddr, + const std::map& 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, stinkytofu::MUBUFScope::SCOPE_NONE); + + if (!mod.offen && !isOffVaddrContainer(vaddr)) { + mod.offen = 1; + } + return mod; } stinkytofu::SMEMModifiers convertSMEMModifiers(const rocisa::SMEMModifiers& rocMod, @@ -553,10 +601,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(inst)) { + stinkyInst->addModifier( + buildMUBUFModifiersForBufferOp(typed->mubuf, typed->vaddr.get(), asmCaps)); + } + else if (auto typed = dynamic_cast(inst)) { + stinkyInst->addModifier( + 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, diff --git a/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp b/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp index 623ad6aa92c..39388f9fdda 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.scope != MUBUFScope::SCOPE_NONE) { + os << ", scope = \"" << toString(mod.scope) << "\""; + } os << " }"; return true; } @@ -377,10 +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") { + 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))); + getBool(fields, "nt", false), getBool(fields, "lds", false), false, + 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 43761f6bd22..c4884623a38 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.scope != MUBUFScope::SCOPE_NONE) { + os << " scope:" << toString(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..8bdcc69facc 100644 --- a/shared/stinkytofu/tests/unit/asm/StinkyAsmEmitterTest.cpp +++ b/shared/stinkytofu/tests/unit/asm/StinkyAsmEmitterTest.cpp @@ -703,3 +703,32 @@ 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, /*scope=*/MUBUFScope::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); +}