From c36ec243142095051b85728450a733c77e0c7292 Mon Sep 17 00:00:00 2001 From: AndySu12 Date: Thu, 4 Jun 2026 07:50:06 -0500 Subject: [PATCH 1/3] Add gfx1250 support for AdaptiveGemm and AdaptiveGemmNTAB --- .../tensilelite/Tensile/KernelWriter.py | 15 +++ .../gemm/gfx12/agntab_coverage_gfx1250.yaml | 105 ++++++++++++++++++ .../tensilelite/rocisa/CMakeLists.txt | 4 +- shared/stinkytofu/CMakeLists.txt | 10 +- .../conversion/rocisa/ToStinkyTofuUtils.cpp | 12 +- 5 files changed, 138 insertions(+), 8 deletions(-) create mode 100644 projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/agntab_coverage_gfx1250.yaml diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py index a63a57da1e7c..2fc0167b7394 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py @@ -5736,6 +5736,14 @@ def _kernelBody(pack, packPre, nta, ntb): originalNta = tensorParametersA["NonTemporal"] originalNtb = tensorParametersB["NonTemporal"] + # gfx1250: NonTemporal only selects cache SCOPE; the actual non-temporal + # access comes from the TemporalHint (TH_NT) field, read per-load from + # kernel["TemporalHint{A,B}"]. Vary it per body alongside NonTemporal. + hasTH = self.states.asmCaps.get("HasTHModifier", False) + if hasTH: + originalThA = kernel["TemporalHintA"] + originalThB = kernel["TemporalHintB"] + ntCombos = [[0, 0], [0, 4], [4, 0]] ntLabels = [Label("LoopBody_NTA{}_NTB{}".format(nta, ntb), "") for nta, ntb in ntCombos] ntLabelDone = Label("LoopBody_NTA_NTB_Done", "") @@ -5799,6 +5807,10 @@ def _restoreNtabState(): module.add(ntLabels[idx]) tensorParametersA["NonTemporal"] = nta tensorParametersB["NonTemporal"] = ntb + if hasTH: + # TH_NT(1) for the NT body, TH_RT(0) otherwise. + kernel["TemporalHintA"] = 1 if nta else 0 + kernel["TemporalHintB"] = 1 if ntb else 0 _kernelBody(pack, packPre, nta, ntb) # All paths but the last one need to skip the rest. Use long # branches everywhere because each kernelBody can easily exceed @@ -5812,6 +5824,9 @@ def _restoreNtabState(): tensorParametersA["NonTemporal"] = originalNta tensorParametersB["NonTemporal"] = originalNtb + if hasTH: + kernel["TemporalHintA"] = originalThA + kernel["TemporalHintB"] = originalThB if kernel["ExpertSchedulingMode"] > 0: module.add(SSetRegIMM32B32(dst=HWRegContainer(reg="26", value=[0,2]), src=0x0, comment="enable hardware dependency checking")) diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/agntab_coverage_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/agntab_coverage_gfx1250.yaml new file mode 100644 index 000000000000..849c3d835cba --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/agntab_coverage_gfx1250.yaml @@ -0,0 +1,105 @@ +TestParameters: + marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx950, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201] # gfx1250 only (StinkyTofu / SIA=4) + +GlobalParameters: + NumElementsToValidate: -1 # full validation: AGNTAB/AG multi-body bugs are silent miscompute + NumWarmups: 1 + MaxEnqueuesPerSync: 1 + SleepPercent: 0 + PrintSolutionRejectionReason: True + DataInitTypeA: 23 + DataInitTypeB: 24 + DataInitTypeBeta: 0 + DataInitTypeAlpha: 1 + CSVExportWinner: 1 + CSVMergeSameProblemID: 1 + Device: 0 + MaxWorkspaceSize: 134217728 + KernelTime: True + RotatingBufferSize: 512 + Architecture: gfx1250 + CodeObjectVersion: 4 + ISA: + - [12, 5, 0] + +.AdaptiveBase: &adaptive_base + CustomKernels: + InitialSolutionParameters: + BenchmarkCommonParameters: + - KernelLanguage: ["Assembly"] + ForkParameters: + - MatrixInstruction: + - [16, 16, 32, 1, 1, 1, 1, 1, 1] # gfx1250 WMMA bf16 16x16x32 + - WavefrontSize: [32] + - DepthU: [32] + - PrefetchGlobalRead: [1] + - PrefetchLocalRead: [1] + - TransposeLDS: [-1] + - WorkGroupMapping: [1] + - 1LDSBuffer: [-1] + - GlobalSplitU: [1] + - GlobalSplitUAlgorithm: ["MultipleBuffer"] + # ============ StinkyTofu scheduling + AGNTAB/AG multi-body under test ============ + - ScheduleIterAlg: [4] # SIA=4 -> StinkyTofu full optimization + - AdaptiveGemm: [0, 1] # 1 emits multiple store blocks (epilogue VW dispatch) + - AdaptiveGemmNTAB: [0, 1] # 1 emits the 3-way main-loop bodies StinkyTofu must schedule + - InitCIterWmma: [0, 1] # 0=off vs 1=force WMMA C-init iter (gfx1250 HasWMMA_AccImmZero=1); crosses AGNTAB initC/openLoop path + BenchmarkJoinParameters: + BenchmarkFinalParameters: + - ProblemSizes: + - Exact: [128, 128, 1, 128] # baseline -> NTA0_NTB0 + - Exact: [32768, 128, 1, 192] # NTA4 path -> NTA4_NTB0 + - Exact: [128, 16384, 1, 256] # NTB4 path -> NTA0_NTB4 + +BenchmarkProblems: + # ================================================================= + # AGNTAB x AG x StinkyTofu (SIA=4) coverage -- 4 transpose x 3 sizes + # Focus: exercise the multi-body main loop + multi-store epilogue through StinkyTofu. + # ================================================================= + - # NN adaptive + - OperationType: GEMM + DataType: b + DestDataType: b + ComputeDataType: s + HighPrecisionAccumulate: True + TransposeA: False + TransposeB: False + UseBeta: True + Batched: True + - <<: *adaptive_base + + - # TN adaptive + - OperationType: GEMM + DataType: b + DestDataType: b + ComputeDataType: s + HighPrecisionAccumulate: True + TransposeA: True + TransposeB: False + UseBeta: True + Batched: True + - <<: *adaptive_base + + - # NT adaptive + - OperationType: GEMM + DataType: b + DestDataType: b + ComputeDataType: s + HighPrecisionAccumulate: True + TransposeA: False + TransposeB: True + UseBeta: True + Batched: True + - <<: *adaptive_base + + - # TT adaptive + - OperationType: GEMM + DataType: b + DestDataType: b + ComputeDataType: s + HighPrecisionAccumulate: True + TransposeA: True + TransposeB: True + UseBeta: True + Batched: True + - <<: *adaptive_base diff --git a/projects/hipblaslt/tensilelite/rocisa/CMakeLists.txt b/projects/hipblaslt/tensilelite/rocisa/CMakeLists.txt index 2b44eb7c293a..335b93bf6758 100644 --- a/projects/hipblaslt/tensilelite/rocisa/CMakeLists.txt +++ b/projects/hipblaslt/tensilelite/rocisa/CMakeLists.txt @@ -6,7 +6,9 @@ if(CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR) cmake_minimum_required(VERSION 3.25) - include(${CMAKE_CURRENT_LIST_DIR}/../../../../cmake/modules/default_amdclang.cmake) + if(EXISTS ${CMAKE_CURRENT_LIST_DIR}/../../../../cmake/modules/default_amdclang.cmake) + include(${CMAKE_CURRENT_LIST_DIR}/../../../../cmake/modules/default_amdclang.cmake) + endif() project(rocisa LANGUAGES CXX) set(ROCISA_STANDALONE ON) diff --git a/shared/stinkytofu/CMakeLists.txt b/shared/stinkytofu/CMakeLists.txt index 76b8fa550cdf..0323f9df4e95 100644 --- a/shared/stinkytofu/CMakeLists.txt +++ b/shared/stinkytofu/CMakeLists.txt @@ -4,7 +4,9 @@ cmake_minimum_required(VERSION 3.16) if(CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR) - include(${CMAKE_CURRENT_LIST_DIR}/../../cmake/modules/default_amdclang.cmake) + if(EXISTS ${CMAKE_CURRENT_LIST_DIR}/../../cmake/modules/default_amdclang.cmake) + include(${CMAKE_CURRENT_LIST_DIR}/../../cmake/modules/default_amdclang.cmake) + endif() endif() set(STINKYTOFU_VERSION "0.1.0") @@ -351,8 +353,10 @@ install(FILES # --------------------------------------------------------------------------- # option(ENABLE_CLANG_TIDY "Enable clang-tidy targets" OFF) list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/../../cmake/modules") -include(ClangTidy) -add_clang_tidy_custom_target() +if(ENABLE_CLANG_TIDY) + include(ClangTidy) + add_clang_tidy_custom_target() +endif() # Build-tree export so find_package() works without installing. # Write targets to stinkytofu-targets.cmake (not stinkytofuConfig.cmake) to diff --git a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp index 1ddf1d5d6aec..f3d5257ca0fb 100644 --- a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp +++ b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp @@ -1095,7 +1095,11 @@ static std::shared_ptr toStinkyTofuModule( }; // Auto-detect the loopWithPrefetch region: from the first global read or - // tensor load item up to and including Module("loopBody"). + // tensor load item up to and including the LAST Module("loopBody"). + // AdaptiveGemmNTAB emits the main loop multiple times (one Module("loopBody") + // per NT body); the region must span all of them, not just the first, so the + // whole multi-body main loop is scheduled. For non-AGNTAB kernels there is + // exactly one "loopBody", so this matches the previous single-body behavior. int pgrStartIdx = -1; int loopBodyIdx = -1; for (int i = 0; i < static_cast(module.itemList.size()); ++i) { @@ -1105,8 +1109,7 @@ static std::shared_ptr toStinkyTofuModule( } if (const auto* subMod = dynamic_cast(item.get())) { if (subMod->name == "loopBody") { - loopBodyIdx = i; - break; + loopBodyIdx = i; // keep updating: take the last (AGNTAB: N bodies) } } } @@ -1172,7 +1175,8 @@ static std::shared_ptr toStinkyTofuModule( static const std::string kScope = "expertScheduleMode2"; // Traverse top-level items, injecting the loopWithPrefetch group name - // for items in the detected prefetch region [pgrStartIdx, loopBodyIdx]. + // for items in the detected prefetch region [pgrStartIdx, loopBodyIdx] + // (spans all AGNTAB main-loop bodies when present). for (int i = 0; i < static_cast(module.itemList.size()); ++i) { const auto& item = module.itemList[i]; const bool inPGR = hasPGR && (i >= pgrStartIdx && i <= loopBodyIdx); From 8b1a24e4a65db12288695046dede2b06d88f1e7d Mon Sep 17 00:00:00 2001 From: AndySu12 Date: Tue, 9 Jun 2026 17:23:28 +0800 Subject: [PATCH 2/3] [hipblaslt] Carry gfx1250 th:/nv modifiers through StinkyTofu emit StinkyTofu's MUBUF modifier model did not carry the gfx1250 TemporalHint (th:) or NonVolatile (nv) fields, so converting a rocisa module via toStinkyTofuModule and rendering it with emitAssembly dropped them while keeping scope:. Because every gfx1250 kernel is emitted through StinkyTofu for any ScheduleIterAlg, this silently disabled the AdaptiveGemmNTAB per-body non-temporal hint (th:TH_LOAD_NT) on all gfx1250 kernels. Add TemporalHint/NonVolatile to StinkyTofu's MUBUFModifiers across the model, the rocisa->StinkyTofu conversion, the asm emitter, and the .stir round-trip, mirroring rocisa's encoding and isStore-aware naming. StinkyTofu emits th:/nv purely from the modifier values; capability gating lives in rocisa (container.hpp) upstream. Add regression tests for the th: and nv paths. --- .../tensilelite/Tensile/KernelWriter.py | 4 - .../gemm/gfx12/agntab_coverage_gfx1250.yaml | 22 ++-- .../tensilelite/rocisa/CMakeLists.txt | 4 +- .../tensilelite/rocisa/test/test_mubuf.py | 112 +++++++++++++++++- shared/stinkytofu/CMakeLists.txt | 10 +- .../stinkytofu/ir/asm/StinkyModifiers.hpp | 17 ++- .../conversion/rocisa/ToStinkyTofuUtils.cpp | 24 ++-- .../serialization/asm/ModifierSerializer.cpp | 11 +- .../serialization/asm/StinkyAsmEmitter.cpp | 3 + 9 files changed, 167 insertions(+), 40 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py index 2fc0167b7394..ed06561a2528 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py @@ -5736,9 +5736,6 @@ def _kernelBody(pack, packPre, nta, ntb): originalNta = tensorParametersA["NonTemporal"] originalNtb = tensorParametersB["NonTemporal"] - # gfx1250: NonTemporal only selects cache SCOPE; the actual non-temporal - # access comes from the TemporalHint (TH_NT) field, read per-load from - # kernel["TemporalHint{A,B}"]. Vary it per body alongside NonTemporal. hasTH = self.states.asmCaps.get("HasTHModifier", False) if hasTH: originalThA = kernel["TemporalHintA"] @@ -5808,7 +5805,6 @@ def _restoreNtabState(): tensorParametersA["NonTemporal"] = nta tensorParametersB["NonTemporal"] = ntb if hasTH: - # TH_NT(1) for the NT body, TH_RT(0) otherwise. kernel["TemporalHintA"] = 1 if nta else 0 kernel["TemporalHintB"] = 1 if ntb else 0 _kernelBody(pack, packPre, nta, ntb) diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/agntab_coverage_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/agntab_coverage_gfx1250.yaml index 849c3d835cba..792347bcabf9 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/agntab_coverage_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/agntab_coverage_gfx1250.yaml @@ -1,8 +1,8 @@ TestParameters: - marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx950, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201] # gfx1250 only (StinkyTofu / SIA=4) + marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx950, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201] GlobalParameters: - NumElementsToValidate: -1 # full validation: AGNTAB/AG multi-body bugs are silent miscompute + NumElementsToValidate: -1 NumWarmups: 1 MaxEnqueuesPerSync: 1 SleepPercent: 0 @@ -11,16 +11,14 @@ GlobalParameters: DataInitTypeB: 24 DataInitTypeBeta: 0 DataInitTypeAlpha: 1 - CSVExportWinner: 1 - CSVMergeSameProblemID: 1 + CSVExportWinner: True + CSVMergeSameProblemID: True Device: 0 MaxWorkspaceSize: 134217728 KernelTime: True RotatingBufferSize: 512 Architecture: gfx1250 - CodeObjectVersion: 4 - ISA: - - [12, 5, 0] + CodeObjectVersion: "4" .AdaptiveBase: &adaptive_base CustomKernels: @@ -40,10 +38,10 @@ GlobalParameters: - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] # ============ StinkyTofu scheduling + AGNTAB/AG multi-body under test ============ - - ScheduleIterAlg: [4] # SIA=4 -> StinkyTofu full optimization - - AdaptiveGemm: [0, 1] # 1 emits multiple store blocks (epilogue VW dispatch) - - AdaptiveGemmNTAB: [0, 1] # 1 emits the 3-way main-loop bodies StinkyTofu must schedule - - InitCIterWmma: [0, 1] # 0=off vs 1=force WMMA C-init iter (gfx1250 HasWMMA_AccImmZero=1); crosses AGNTAB initC/openLoop path + - ScheduleIterAlg: [0, 1, 2, 3, 4] + - AdaptiveGemm: [0, 1] + - AdaptiveGemmNTAB: [0, 1] + - InitCIterWmma: [0, 1] BenchmarkJoinParameters: BenchmarkFinalParameters: - ProblemSizes: @@ -53,7 +51,7 @@ GlobalParameters: BenchmarkProblems: # ================================================================= - # AGNTAB x AG x StinkyTofu (SIA=4) coverage -- 4 transpose x 3 sizes + # AGNTAB x AG x StinkyTofu coverage -- 4 transpose x 3 sizes # Focus: exercise the multi-body main loop + multi-store epilogue through StinkyTofu. # ================================================================= - # NN adaptive diff --git a/projects/hipblaslt/tensilelite/rocisa/CMakeLists.txt b/projects/hipblaslt/tensilelite/rocisa/CMakeLists.txt index 335b93bf6758..2b44eb7c293a 100644 --- a/projects/hipblaslt/tensilelite/rocisa/CMakeLists.txt +++ b/projects/hipblaslt/tensilelite/rocisa/CMakeLists.txt @@ -6,9 +6,7 @@ if(CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR) cmake_minimum_required(VERSION 3.25) - if(EXISTS ${CMAKE_CURRENT_LIST_DIR}/../../../../cmake/modules/default_amdclang.cmake) - include(${CMAKE_CURRENT_LIST_DIR}/../../../../cmake/modules/default_amdclang.cmake) - endif() + include(${CMAKE_CURRENT_LIST_DIR}/../../../../cmake/modules/default_amdclang.cmake) project(rocisa LANGUAGES CXX) set(ROCISA_STANDALONE ON) diff --git a/projects/hipblaslt/tensilelite/rocisa/test/test_mubuf.py b/projects/hipblaslt/tensilelite/rocisa/test/test_mubuf.py index 9d60add2a148..5abbc5c5f5ad 100644 --- a/projects/hipblaslt/tensilelite/rocisa/test/test_mubuf.py +++ b/projects/hipblaslt/tensilelite/rocisa/test/test_mubuf.py @@ -28,7 +28,7 @@ import rocisa from rocisa.code import Module, SignatureBase from rocisa.container import MUBUFModifiers, sgpr, vgpr -from rocisa.enum import CacheScope +from rocisa.enum import CacheScope, TemporalHint, NonVolatile from rocisa.instruction import ( BufferAtomicAddF32, BufferLoadB32, @@ -156,6 +156,116 @@ def test_mubuf_scope_modifiers_stinkytofu(_mubuf_scope_asm): ) +@pytest.fixture(scope="module") +def _mubuf_th_asm() -> str: + mod = Module("mubuf_th_modifiers") + mod.add( + BufferLoadB32( + dst=vgpr(13), + vaddr=vgpr(33), + saddr=sgpr(64, 4), + soffset=sgpr(47), + mubuf=MUBUFModifiers(offen=True, scope=CacheScope.SCOPE_CU, th=TemporalHint.TH_NT), + ) + ) + mod.add( + BufferStoreB32( + src=vgpr(12), + vaddr=vgpr(32), + saddr=sgpr(60, 4), + soffset=sgpr(46), + mubuf=MUBUFModifiers(offen=True, isStore=True, scope=CacheScope.SCOPE_CU, th=TemporalHint.TH_NT), + ) + ) + mod.setParent() + + sig = SignatureBase( + kernelName="mubuf_th_modifiers", + kernArgsVersion=1, + codeObjectVersion="4", + groupSegmentSize=0, + sgprWorkGroup=(1, 1, 0), + vgprWorkItem=0, + flatWorkGroupSize=64, + numSgprPreload=0, + ) + + st = rocisa.toStinkyTofuModule( + mod, _ISA, "mubuf_th_modifiers", signature=sig, options={"OptLevel": 0} + ) + st.runOptimizationPipeline() + return st.emitAssembly() + + +def test_mubuf_temporal_hint_modifiers_stinkytofu(_mubuf_th_asm): + assert re.search( + r"buffer_load_b32 v13, v33, s\[64:67\], s47 offen offset:0 scope:SCOPE_CU th:TH_LOAD_NT", + _mubuf_th_asm, + ) + assert re.search( + r"buffer_store_b32 v12, v32, s\[60:63\], s46 offen offset:0 scope:SCOPE_CU th:TH_STORE_NT", + _mubuf_th_asm, + ) + + +@pytest.fixture(scope="module") +def _mubuf_nv_asm() -> str: + mod = Module("mubuf_nv_modifiers") + mod.add( + BufferLoadB32( + dst=vgpr(13), + vaddr=vgpr(33), + saddr=sgpr(64, 4), + soffset=sgpr(47), + mubuf=MUBUFModifiers(offen=True, scope=CacheScope.SCOPE_CU, nv=NonVolatile.NV), + ) + ) + mod.add( + BufferStoreB32( + src=vgpr(12), + vaddr=vgpr(32), + saddr=sgpr(60, 4), + soffset=sgpr(46), + mubuf=MUBUFModifiers( + offen=True, + isStore=True, + scope=CacheScope.SCOPE_CU, + th=TemporalHint.TH_NT, + nv=NonVolatile.NV, + ), + ) + ) + mod.setParent() + + sig = SignatureBase( + kernelName="mubuf_nv_modifiers", + kernArgsVersion=1, + codeObjectVersion="4", + groupSegmentSize=0, + sgprWorkGroup=(1, 1, 0), + vgprWorkItem=0, + flatWorkGroupSize=64, + numSgprPreload=0, + ) + + st = rocisa.toStinkyTofuModule( + mod, _ISA, "mubuf_nv_modifiers", signature=sig, options={"OptLevel": 0} + ) + st.runOptimizationPipeline() + return st.emitAssembly() + + +def test_mubuf_non_volatile_modifiers_stinkytofu(_mubuf_nv_asm): + assert re.search( + r"buffer_load_b32 v13, v33, s\[64:67\], s47 offen offset:0 scope:SCOPE_CU nv", + _mubuf_nv_asm, + ) + assert re.search( + r"buffer_store_b32 v12, v32, s\[60:63\], s46 offen offset:0 scope:SCOPE_CU th:TH_STORE_NT nv", + _mubuf_nv_asm, + ) + + def test_rocisa_load_null_soffset_adds_offen(): inst = BufferLoadB32( dst=vgpr(13), diff --git a/shared/stinkytofu/CMakeLists.txt b/shared/stinkytofu/CMakeLists.txt index 0323f9df4e95..76b8fa550cdf 100644 --- a/shared/stinkytofu/CMakeLists.txt +++ b/shared/stinkytofu/CMakeLists.txt @@ -4,9 +4,7 @@ cmake_minimum_required(VERSION 3.16) if(CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR) - if(EXISTS ${CMAKE_CURRENT_LIST_DIR}/../../cmake/modules/default_amdclang.cmake) - include(${CMAKE_CURRENT_LIST_DIR}/../../cmake/modules/default_amdclang.cmake) - endif() + include(${CMAKE_CURRENT_LIST_DIR}/../../cmake/modules/default_amdclang.cmake) endif() set(STINKYTOFU_VERSION "0.1.0") @@ -353,10 +351,8 @@ install(FILES # --------------------------------------------------------------------------- # option(ENABLE_CLANG_TIDY "Enable clang-tidy targets" OFF) list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/../../cmake/modules") -if(ENABLE_CLANG_TIDY) - include(ClangTidy) - add_clang_tidy_custom_target() -endif() +include(ClangTidy) +add_clang_tidy_custom_target() # Build-tree export so find_package() works without installing. # Write targets to stinkytofu-targets.cmake (not stinkytofuConfig.cmake) to diff --git a/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp b/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp index d49c0b23ed9a..e5a6a391ec2d 100644 --- a/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp +++ b/shared/stinkytofu/include/stinkytofu/ir/asm/StinkyModifiers.hpp @@ -177,6 +177,17 @@ inline TemporalHint parseTemporalHint(std::string_view th) { return TemporalHint::TH_NONE; } +enum class NonVolatile : uint8_t { NV_NONE = 0, NV = 1 }; + +inline std::string_view toString(NonVolatile nv) { + return nv == NonVolatile::NV ? "nv" : ""; +} + +// Inverse of toString(): assembly token -> enum. +inline NonVolatile parseNonVolatile(std::string_view nv) { + return nv == "nv" ? NonVolatile::NV : NonVolatile::NV_NONE; +} + // 9-bit DPP permutation control selector (matches the hardware dpp_ctrl field). // Three encoding shapes: // singleton — the named value IS the encoding (e.g. ROW_MIRROR = 0x140) @@ -404,7 +415,7 @@ struct MUBUFModifiers : public TypedModifier { bool nt = false, bool lds = false, bool isStore = false, bool hasMUBUFConst = false, bool hasGLCModifier = false, bool hasSC0Modifier = false, MUBUFScope scope = MUBUFScope::SCOPE_NONE, - TemporalHint th = TemporalHint::TH_NONE) + TemporalHint th = TemporalHint::TH_NONE, NonVolatile nv = NonVolatile::NV_NONE) : TypedModifier(), offset12(offset12), offen(offen), @@ -417,7 +428,8 @@ struct MUBUFModifiers : public TypedModifier { hasGLCModifier(hasGLCModifier), hasSC0Modifier(hasSC0Modifier), scope(scope), - th(th) {} + th(th), + nv(nv) {} int offset12; uint32_t offen : 1; @@ -431,6 +443,7 @@ struct MUBUFModifiers : public TypedModifier { uint32_t hasSC0Modifier : 1; MUBUFScope scope; TemporalHint th; + NonVolatile nv; }; // Carries just the cache scope token for SOPP-format memory fences such as diff --git a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp index f3d5257ca0fb..f3bd22cbecf0 100644 --- a/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp +++ b/shared/stinkytofu/src/conversion/rocisa/ToStinkyTofuUtils.cpp @@ -130,6 +130,17 @@ stinkytofu::TemporalHint convertTemporalHint(rocisa::TemporalHint th) { } } +stinkytofu::NonVolatile convertNonVolatile(rocisa::NonVolatile nv) { + switch (nv) { + case rocisa::NonVolatile::NV_NONE: + return stinkytofu::NonVolatile::NV_NONE; + case rocisa::NonVolatile::NV: + return stinkytofu::NonVolatile::NV; + default: + return stinkytofu::NonVolatile::NV_NONE; + } +} + stinkytofu::MUBUFModifiers convertMUBUFModifiers(const rocisa::MUBUFModifiers& rocMod, const std::map& asmCaps) { bool hasMUBUFConst = asmCaps.count("HasMUBUFConst") && asmCaps.at("HasMUBUFConst"); @@ -137,9 +148,10 @@ stinkytofu::MUBUFModifiers convertMUBUFModifiers(const rocisa::MUBUFModifiers& r bool hasSC0Modifier = asmCaps.count("HasSC0Modifier") && asmCaps.at("HasSC0Modifier"); stinkytofu::MUBUFScope scope = convertMUBUFScope(rocMod.scope); stinkytofu::TemporalHint th = convertTemporalHint(rocMod.th); + stinkytofu::NonVolatile nv = convertNonVolatile(rocMod.nv); return stinkytofu::MUBUFModifiers(rocMod.offen, rocMod.offset12, rocMod.glc, rocMod.slc, rocMod.nt, rocMod.lds, rocMod.isStore, hasMUBUFConst, - hasGLCModifier, hasSC0Modifier, scope, th); + hasGLCModifier, hasSC0Modifier, scope, th, nv); } /// Returns true when vaddr is the MUBUF "off" keyword. @@ -1095,11 +1107,7 @@ static std::shared_ptr toStinkyTofuModule( }; // Auto-detect the loopWithPrefetch region: from the first global read or - // tensor load item up to and including the LAST Module("loopBody"). - // AdaptiveGemmNTAB emits the main loop multiple times (one Module("loopBody") - // per NT body); the region must span all of them, not just the first, so the - // whole multi-body main loop is scheduled. For non-AGNTAB kernels there is - // exactly one "loopBody", so this matches the previous single-body behavior. + // tensor load item up to and including Module("loopBody"). int pgrStartIdx = -1; int loopBodyIdx = -1; for (int i = 0; i < static_cast(module.itemList.size()); ++i) { @@ -1109,7 +1117,7 @@ static std::shared_ptr toStinkyTofuModule( } if (const auto* subMod = dynamic_cast(item.get())) { if (subMod->name == "loopBody") { - loopBodyIdx = i; // keep updating: take the last (AGNTAB: N bodies) + loopBodyIdx = i; // keep updating loop bodies } } } @@ -1176,7 +1184,7 @@ static std::shared_ptr toStinkyTofuModule( // Traverse top-level items, injecting the loopWithPrefetch group name // for items in the detected prefetch region [pgrStartIdx, loopBodyIdx] - // (spans all AGNTAB main-loop bodies when present). + // (spans all main-loop bodies when present). for (int i = 0; i < static_cast(module.itemList.size()); ++i) { const auto& item = module.itemList[i]; const bool inPGR = hasPGR && (i >= pgrStartIdx && i <= loopBodyIdx); diff --git a/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp b/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp index 25fa52aee6eb..efc1c1b70a8c 100644 --- a/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp +++ b/shared/stinkytofu/src/serialization/asm/ModifierSerializer.cpp @@ -173,7 +173,10 @@ bool serializeVisit(const MUBUFModifiers& mod, std::ostream& os) { os << ", scope = \"" << toString(mod.scope) << "\""; } if (hasTemporalHint(mod.th)) { - os << ", th = \"" << toString(mod.th) << "\""; + os << ", th = \"" << toString(mod.th, mod.isStore) << "\""; + } + if (mod.nv != NonVolatile::NV_NONE) { + os << ", nv = \"" << toString(mod.nv) << "\""; } os << " }"; return true; @@ -454,11 +457,13 @@ void deserializeVisit(StinkyInstruction* inst, const std::string& attrKey, } else if (attrKey == "mod.mubuf") { MUBUFScope scope = parseMUBUFScope(getStr(fields, "scope", "")); TemporalHint th = parseTemporalHint(getStr(fields, "th", "")); + NonVolatile nv = parseNonVolatile(getStr(fields, "nv", "")); + bool isStore = getStr(fields, "th", "").rfind("TH_STORE_", 0) == 0; 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, th)); + getBool(fields, "nt", false), getBool(fields, "lds", false), isStore, + false, false, false, scope, th, nv)); } else if (attrKey == "mod.cache_scope") { inst->addModifier(CacheScopeModifiers(parseMUBUFScope(getStr(fields, "scope", "")))); } else if (attrKey == "mod.smem") { diff --git a/shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp b/shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp index 4ce311abb188..25963337d0c9 100644 --- a/shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp +++ b/shared/stinkytofu/src/serialization/asm/StinkyAsmEmitter.cpp @@ -223,6 +223,9 @@ inline std::ostream& operator<<(std::ostream& os, const MUBUFModifiers& mubufMod } else if (mubufMod.nt) { os << " nt"; } + if (mubufMod.nv != NonVolatile::NV_NONE) { + os << " " << toString(mubufMod.nv); + } if (mubufMod.lds) { os << " lds"; } From 039fe23759ed2ef875ac4d06118feae1539c0625 Mon Sep 17 00:00:00 2001 From: AndySu12 Date: Wed, 17 Jun 2026 18:43:37 +0800 Subject: [PATCH 3/3] [hipblaslt] Switch per-body gfx1250 TemporalHint to tensorParameters working copy AdaptiveGemmNTAB emits multiple loop bodies whose A/B global loads differ only by their non-temporal/temporal-hint state. NonTemporal already switches this per body via the transient tensorParameters["NonTemporal"] working copy; TemporalHint, however, was toggled by mutating the kernel solution parameters kernel["TemporalHintA"/"TemporalHintB"] directly, which is an architectural violation (solution parameters are immutable inputs, not per-body scratch). Mirror the NonTemporal pattern: seed tP["TemporalHint"] alongside tP["NonTemporal"] at tensorParameters init, read it on the A/B global-read paths (TemporalHint(tP["TemporalHint"])), and snapshot/set/restore it at the tensorParameters level inside the NTAB body loop. The kernel-level TemporalHint solution parameters are no longer mutated. Behavior is unchanged: the same per-body int values flow through the same TemporalHint enum into decodeNonTemporal; only the storage location moves from the kernel dict to the per-tensor working copy. The kernel-direct load path (non-tP) keeps reading the solution parameter as before. --- .../hipblaslt/tensilelite/Tensile/KernelWriter.py | 13 +++++++------ .../tensilelite/Tensile/KernelWriterAssembly.py | 4 ++-- 2 files changed, 9 insertions(+), 8 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py index ed06561a2528..7dad49ea7c26 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py @@ -5738,8 +5738,8 @@ def _kernelBody(pack, packPre, nta, ntb): hasTH = self.states.asmCaps.get("HasTHModifier", False) if hasTH: - originalThA = kernel["TemporalHintA"] - originalThB = kernel["TemporalHintB"] + originalThA = tensorParametersA["TemporalHint"] + originalThB = tensorParametersB["TemporalHint"] ntCombos = [[0, 0], [0, 4], [4, 0]] ntLabels = [Label("LoopBody_NTA{}_NTB{}".format(nta, ntb), "") for nta, ntb in ntCombos] @@ -5805,8 +5805,8 @@ def _restoreNtabState(): tensorParametersA["NonTemporal"] = nta tensorParametersB["NonTemporal"] = ntb if hasTH: - kernel["TemporalHintA"] = 1 if nta else 0 - kernel["TemporalHintB"] = 1 if ntb else 0 + tensorParametersA["TemporalHint"] = 1 if nta else 0 + tensorParametersB["TemporalHint"] = 1 if ntb else 0 _kernelBody(pack, packPre, nta, ntb) # All paths but the last one need to skip the rest. Use long # branches everywhere because each kernelBody can easily exceed @@ -5821,8 +5821,8 @@ def _restoreNtabState(): tensorParametersA["NonTemporal"] = originalNta tensorParametersB["NonTemporal"] = originalNtb if hasTH: - kernel["TemporalHintA"] = originalThA - kernel["TemporalHintB"] = originalThB + tensorParametersA["TemporalHint"] = originalThA + tensorParametersB["TemporalHint"] = originalThB if kernel["ExpertSchedulingMode"] > 0: module.add(SSetRegIMM32B32(dst=HWRegContainer(reg="26", value=[0,2]), src=0x0, comment="enable hardware dependency checking")) @@ -9932,6 +9932,7 @@ def getTensorParameters(self, tP, kernel, itP, cM): tP["wtc"] = itP[cM].writeTileDimComponents # write vector components along tile dimension tP["idx"] = kernel["ProblemType"]["Index%d"%tP["tensorIdx"]] # index 0 is tile dimension belonging to A. Note 'idx' may not be in tP['ia']. tP["NonTemporal"] = kernel["NonTemporal%s"%cM] # non-temporal read type + tP["TemporalHint"] = kernel.get("TemporalHint%s"%cM, 0) # temporal-hint read type tP["shiftGR"] = 0 if (tP["bpeGR"] >= tP["bpeDS"]) else int(tP["glvw"] // 2 * (tP["bpeDS"] / self.states.bpr)) # Shift global read register for cvt spaces tP["bpeRatio"] = tP["bpeDS"] // tP["bpeGR"] if tP["bpeGR"] < tP["bpeDS"] else 1 # g2lIdx multiplier diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py index eb25ac4b68e4..4b09e8d1c428 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py @@ -10341,7 +10341,7 @@ def globalReadGuardKBody(tP, optParams = None): dataType = kernel["ProblemType"]["MacDataType%s"%tcDataType if (tcDataType=='A' or tcDataType=='B') else "DataType"] isGlc, isSlc, isNT, scope, th, nv = decodeNonTemporal( - self.states.asmCaps, tP["NonTemporal"], _temporalHint(kernel, tc), _nonVolatile(kernel, tc)) + self.states.asmCaps, tP["NonTemporal"], TemporalHint(tP["TemporalHint"]), _nonVolatile(kernel, tc)) isLds = True if (kernel["DirectToLds%s"%tc] and not kernel["NonDTLTailLoop%s"%tc]) else False isTr = (tc == "A" or tc == "B") and kernel["enableGLTr%s"%tc] is16b = dataType.isHalf() or dataType.isBFloat16() @@ -11276,7 +11276,7 @@ def globalReadBody(tP): bpe = tP["bpeGR"] if isAB else tP["bpe"] bpl = bpe * tP["glvw"] # bytes per load isGlc, isSlc, isNT, scope, th, nv = decodeNonTemporal( - self.states.asmCaps, tP["NonTemporal"], _temporalHint(kernel, tc), _nonVolatile(kernel, tc)) + self.states.asmCaps, tP["NonTemporal"], TemporalHint(tP["TemporalHint"]), _nonVolatile(kernel, tc)) isLds = True if kernel["DirectToLds%s"%tc] else False isTr = (tc == "A" or tc == "B") and kernel["enableGLTr%s"%tc]