ROCm · jimbochi · Apr 23, 2026 · Apr 24, 2026 · Apr 27, 2026 · Apr 27, 2026
@@ -6651,6 +6651,18 @@ def _restoreNtabState():
       t1a_end = time.perf_counter()
       print2(f"StinkyTofu (1a) toStinkyTofuModule: {t1a_end - t1a_start:.4f}s")
 
+      # --- DEBUG: dump StinkyAsmModule BEFORE optimization pipeline ---
+      # Compares the two paths' logical→asm lowering output before shared passes run.
+      if os.environ.get("DUMP_STINKY_MODULE"):
+        _dump_dir = os.environ.get("DUMP_STINKY_MODULE", "/tmp/stinky_dump")
+        os.makedirs(_dump_dir, exist_ok=True)
+        _backend_tag = os.environ.get("ROCISA_BACKEND", "native")
+        _dump_asm = stModule.emitAssembly()
+        _dump_path = os.path.join(_dump_dir, f"stmodule_{_backend_tag}.s")
+        with open(_dump_path, "w") as _df:
+          _df.write(_dump_asm)
+        print2(f"StinkyTofu DEBUG: dumped pre-pipeline module ({len(_dump_asm)} chars) -> {_dump_path}")
+
       # Run pipeline — builder handles O0 internally (skips optimization,
       # still runs required passes like InsertVgprMsb)
       t1b_start = time.perf_counter()

@@ -161,35 +161,35 @@ __device__ char
     llvm_amdgcn_raw_buffer_load_i8(int32x4_t buffer_resource,
                                    uint32_t  voffset,
                                    uint32_t  soffset,
-                                   int32_t   cache_op) __asm("llvm.amdgcn.raw.buffer.load.i8");
+                                   int32_t   cache_op) __asm("llvm.amdgcn.raw.buffer.load.i8.v4i32");
 
 // 2 bytes
 __device__ float16_t
     llvm_amdgcn_raw_buffer_load_f16(int32x4_t buffer_resource,
                                     uint32_t  voffset,
                                     uint32_t  soffset,
-                                    int32_t   cache_op) __asm("llvm.amdgcn.raw.buffer.load.f16");
+                                    int32_t   cache_op) __asm("llvm.amdgcn.raw.buffer.load.f16.v4i32");
 
 // 4 bytes
 __device__ float32_t
     llvm_amdgcn_raw_buffer_load_f32(int32x4_t buffer_resource,
                                     uint32_t  voffset,
                                     uint32_t  soffset,
-                                    int32_t   cache_op) __asm("llvm.amdgcn.raw.buffer.load.f32");
+                                    int32_t   cache_op) __asm("llvm.amdgcn.raw.buffer.load.f32.v4i32");
 
 // 8 bytes
 __device__ float32x2_t
     llvm_amdgcn_raw_buffer_load_f32x2(int32x4_t buffer_resource,
                                       uint32_t  voffset,
                                       uint32_t  soffset,
-                                      int32_t cache_op) __asm("llvm.amdgcn.raw.buffer.load.v2f32");
+                                      int32_t cache_op) __asm("llvm.amdgcn.raw.buffer.load.v2f32.v4i32");
 
 // 16 bytes
 __device__ float32x4_t
     llvm_amdgcn_raw_buffer_load_f32x4(int32x4_t buffer_resource,
                                       uint32_t  voffset,
                                       uint32_t  soffset,
-                                      int32_t cache_op) __asm("llvm.amdgcn.raw.buffer.load.v4f32");
+                                      int32_t cache_op) __asm("llvm.amdgcn.raw.buffer.load.v4f32.v4i32");
 
 // 4 bytes
 __device__ int32_t
@@ -219,39 +219,39 @@ __device__ void
                                     int32x4_t buffer_resource,
                                     uint32_t  voffset,
                                     uint32_t  soffset,
-                                    int32_t   cache_op) __asm("llvm.amdgcn.raw.buffer.store.i8");
+                                    int32_t   cache_op) __asm("llvm.amdgcn.raw.buffer.store.i8.v4i32");
 
 // 2 bytes
 __device__ void
     llvm_amdgcn_raw_buffer_store_f16(float16_t data,
                                      int32x4_t buffer_resource,
                                      uint32_t  voffset,
                                      uint32_t  soffset,
-                                     int32_t   cache_op) __asm("llvm.amdgcn.raw.buffer.store.f16");
+                                     int32_t   cache_op) __asm("llvm.amdgcn.raw.buffer.store.f16.v4i32");
 
 // 4 bytes
 __device__ void
     llvm_amdgcn_raw_buffer_store_f32(float32_t data,
                                      int32x4_t buffer_resource,
                                      uint32_t  voffset,
                                      uint32_t  soffset,
-                                     int32_t   cache_op) __asm("llvm.amdgcn.raw.buffer.store.f32");
+                                     int32_t   cache_op) __asm("llvm.amdgcn.raw.buffer.store.f32.v4i32");
 
 // 8 bytes
 __device__ void llvm_amdgcn_raw_buffer_store_f32x2(
     float32x2_t data,
     int32x4_t   buffer_resource,
     uint32_t    voffset,
     uint32_t    soffset,
-    int32_t     cache_op) __asm("llvm.amdgcn.raw.buffer.store.v2f32");
+    int32_t     cache_op) __asm("llvm.amdgcn.raw.buffer.store.v2f32.v4i32");
 
 // 16 bytes
 __device__ void llvm_amdgcn_raw_buffer_store_f32x4(
     float32x4_t data,
     int32x4_t   buffer_resource,
     uint32_t    voffset,
     uint32_t    soffset,
-    int32_t     cache_op) __asm("llvm.amdgcn.raw.buffer.store.v4f32");
+    int32_t     cache_op) __asm("llvm.amdgcn.raw.buffer.store.v4f32.v4i32");
 
 ////////////////////////////////////////////////////////////////////////////////////////////////////
 

@@ -133,13 +133,21 @@ if(HIPBLASLT_BUNDLE_PYTHON_DEPS OR ROCISA_STANDALONE)
     if(NOT TARGET stinkytofu::stinkytofu)
         find_package(stinkytofu QUIET)
         if(NOT stinkytofu_FOUND)
-            set(_stinkytofu_sibling "${CMAKE_CURRENT_SOURCE_DIR}/../../../../shared/stinkytofu")
-            if(IS_DIRECTORY "${_stinkytofu_sibling}")
+            set(STINKYTOFU_DIR "${CMAKE_CURRENT_SOURCE_DIR}/../../../../shared/stinkytofu")
+            if(EXISTS "${STINKYTOFU_DIR}/CMakeLists.txt")
+                set(BUILD_SHARED_LIBS ON)
                 set(STINKYTOFU_ENABLE_WERROR ON)
-                add_subdirectory("${_stinkytofu_sibling}" stinkytofu EXCLUDE_FROM_ALL)
+
+                # Force StinkyTofu Python bindings ON because the rocisa->stinkytofu 
+                # adapter relies on the _stinkytofu.so nanobind module.
+                set(STINKYTOFU_BUILD_PYTHON ON CACHE BOOL
+                    "Build StinkyTofu Python bindings (required by ROCISA_BACKEND=stinkytofu adapter)"
+                    FORCE)
+
+                add_subdirectory("${STINKYTOFU_DIR}" stinkytofu EXCLUDE_FROM_ALL)
                 set(_stinkytofu_from_subdirectory TRUE)
             else()
-                message(FATAL_ERROR "stinkytofu not found via find_package or sibling directory")
+                message(FATAL_ERROR "stinkytofu not found via find_package or sibling directory: ${STINKYTOFU_DIR}")
             endif()
         endif()
     endif()
@@ -235,6 +243,52 @@ if(HIPBLASLT_BUNDLE_PYTHON_DEPS OR ROCISA_STANDALONE)
         install(FILES "${CMAKE_CURRENT_BINARY_DIR}/rocisa/_build_info.py" DESTINATION rocisa)
     endif()
 
+    # ---------------------------------------------------------------------
+    # Sibling staging for the stinkytofu Python package.
+    # ---------------------------------------------------------------------
+    if(TARGET stinkytofu_python)
+        add_dependencies(_rocisa stinkytofu_python)
+
+        set_target_properties(stinkytofu_python PROPERTIES
+            LIBRARY_OUTPUT_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/stinkytofu"
+            BUILD_RPATH "$ORIGIN"
+            INSTALL_RPATH "$ORIGIN"
+        )
+
+        set(STINKYTOFU_TOP_DIR "${STINKYTOFU_DIR}")
+        configure_file(
+            "${STINKYTOFU_DIR}/python_module/stinkytofu/__init__.py"
+            "${CMAKE_CURRENT_BINARY_DIR}/stinkytofu/__init__.py"
+            COPYONLY
+        )
+        configure_file(
+            "${STINKYTOFU_DIR}/python_module/stinkytofu/_build_info.py.in"
+            "${CMAKE_CURRENT_BINARY_DIR}/stinkytofu/_build_info.py"
+            @ONLY
+        )
+
+        # -----------------------------------------------------------------
+        # intrinsics.st.bc — runtime bitcode loaded by IntrinsicRegistry
+        # -----------------------------------------------------------------
+        if(TARGET intrinsics_compiled)
+            add_dependencies(_rocisa intrinsics_compiled)
+
+            set(_intrinsics_bc_src "${CMAKE_BINARY_DIR}/intrinsics.st.bc")
+            set(_intrinsics_bc_dst "${CMAKE_CURRENT_BINARY_DIR}/stinkytofu/intrinsics.st.bc")
+            add_custom_command(
+                OUTPUT "${_intrinsics_bc_dst}"
+                COMMAND ${CMAKE_COMMAND} -E copy_if_different
+                        "${_intrinsics_bc_src}" "${_intrinsics_bc_dst}"
+                DEPENDS intrinsics_compiled "${_intrinsics_bc_src}"
+                COMMENT "Staging intrinsics.st.bc next to _stinkytofu.so"
+                VERBATIM
+            )
+            add_custom_target(rocisa_stage_intrinsics ALL
+                DEPENDS "${_intrinsics_bc_dst}")
+            add_dependencies(_rocisa rocisa_stage_intrinsics)
+        endif()
+    endif()
+
     if(ROCISA_STANDALONE)
         # scikit-build-core wheel install rules: place _rocisa and __init__.py into the
         # rocisa/ package directory in the wheel.

@@ -1,6 +1,7 @@
 # Copyright © Advanced Micro Devices, Inc., or its affiliates.
 # SPDX-License-Identifier: MIT
 
+import os
 import sys
 import types
 
@@ -14,26 +15,83 @@
     )
 del _Path
 
-from ._rocisa import *
-from . import _rocisa
-
-# Register nanobind submodules under the rocisa.* namespace so that
-# `from rocisa.enum import X` and `import rocisa.instruction as ri` work.
-for _name, _obj in vars(_rocisa).items():
-    if isinstance(_obj, types.ModuleType) and not _name.startswith("_"):
-        sys.modules.setdefault(f"rocisa.{_name}", _obj)
-
-# Staleness check: only active in source builds.
-# Pre-built packages (wheels, apt) lack _build_info.py — the import is
-# silently skipped. Catching ImportError (not just ModuleNotFoundError)
-# because Python 3.10 raises ImportError for missing relative submodules.
-# The intentional staleness ImportError is raised outside the try/except
-# so it is never swallowed.
-_bi = None
-try:
-    from . import _build_info as _bi
-except ImportError:
-    pass  # Pre-built package — no source tree, skip check
+# ---------------------------------------------------------------------------
+# Backend dispatch
+# ---------------------------------------------------------------------------
+# ``ROCISA_BACKEND=stinkytofu`` redirects ``import rocisa`` to the
+# ``rocisa_stinkytofu_adaptor`` shim (a rocisa-shaped facade backed by the
+# stinkytofu Python binding ``_stinkytofu.so``). Anything else (or unset)
+# keeps the original nanobind bindings in ``_rocisa``.
+
+_BACKEND = os.environ.get("ROCISA_BACKEND", "").strip().lower()
+
+_ADAPTER_PKG = "rocisa_stinkytofu_adaptor"
+
+
+def _load_stinkytofu_adapter() -> bool:
+    """Try to install the rocisa_stinkytofu_adaptor as the ``rocisa`` module.
+
+    Returns True iff we successfully rewired sys.modules; on any failure
+    we fall back to the nanobind bindings silently (the caller decides
+    whether that is acceptable).
+    """
+
+    # Locate ``<repo_root>/projects/hipblaslt/tensilelite/rocisa_stinkytofu_adaptor``
+    # by walking up from this file until we find an ancestor whose basename
+    # is ``projects``; the adapter then lives at
+    # ``<that_parent>/projects/hipblaslt/tensilelite/<_ADAPTER_PKG>/``.
+    #
+    # The adapter is a sibling of ``tensilelite/rocisa/`` on purpose — it is
+    # a *consumer* of the stinkytofu Python binding (a tensilelite-internal
+    # alternative backend), not a piece of ``shared/stinkytofu/`` itself.
+    # Sibling-of-rocisa keeps ``ROCISA_BACKEND=stinkytofu`` purely a
+    # tensilelite concern.
+    #
+    # Works for both the source tree (``<repo>/projects/hipblaslt/
+    # tensilelite/rocisa/rocisa/__init__.py``) and CMake-staged copies
+    # under ``<repo>/projects/hipblaslt/tensilelite/<build_dir>/tensilelite/
+    # rocisa/rocisa/__init__.py`` because in either case walking up from
+    # ``__file__`` eventually hits the ``projects`` directory.
+    repo_root = None
+    cur = os.path.dirname(os.path.abspath(__file__))
+    adapter_rel = os.path.join("projects", "hipblaslt", "tensilelite", _ADAPTER_PKG)
+    while True:
+        parent = os.path.dirname(cur)
+        if parent == cur:  # reached filesystem root
+            break
+        if os.path.basename(cur) == "projects":
+            candidate = parent
+            if os.path.isdir(os.path.join(candidate, adapter_rel)):
+                repo_root = candidate
+                break
+        cur = parent
+    if repo_root is None:
+        return False
+    adapter_parent = os.path.join(repo_root, adapter_rel)
+
+    if not os.path.isdir(os.path.join(adapter_parent, _ADAPTER_PKG)):
+        return False
+
+    if adapter_parent not in sys.path:
+        sys.path.insert(0, adapter_parent)
+
+    try:
+        import rocisa_stinkytofu_adaptor as _adapter  # noqa: F401
+    except Exception:
+        return False
+
+    # Install the adapter as ``rocisa`` and re-export each
+    # ``rocisa_stinkytofu_adaptor.*`` submodule under ``rocisa.*`` in
+    # ``sys.modules``.
+    sys.modules["rocisa"] = _adapter
+    _prefix = f"{_ADAPTER_PKG}."
+    for _name, _obj in vars(_adapter).items():
+        if isinstance(_obj, types.ModuleType) and _obj.__name__.startswith(_prefix):
+            short = _obj.__name__[len(_prefix):]
+            sys.modules[f"rocisa.{short}"] = _obj
+
+    return True
+
 
 def _find_stale_sources(so_path, source_roots, build_dir):
     """Return source files newer than so_path, excluding files under build_dir.
@@ -54,37 +112,70 @@ def _find_stale_sources(so_path, source_roots, build_dir):
     return stale
 
 
-if _bi is not None:
-    from pathlib import Path
+if _BACKEND == "stinkytofu" and _load_stinkytofu_adapter():
+    # stinkytofu adapter active; wiring done inside _load_stinkytofu_adapter.
+    pass
+else:
+    # Default path: original nanobind bindings.
+    from ._rocisa import *  # noqa: F401,F403
+    from . import _rocisa
+
+    # Register nanobind submodules under the rocisa.* namespace so that
+    # `from rocisa.enum import X` and `import rocisa.instruction as ri` work.
+    for _name, _obj in vars(_rocisa).items():
+        if isinstance(_obj, types.ModuleType) and not _name.startswith("_"):
+            sys.modules.setdefault(f"rocisa.{_name}", _obj)
+
+    # Staleness check: only active in source builds.
+    # Pre-built packages (wheels, apt) lack _build_info.py — the import is
+    # silently skipped. Catching ImportError (not just ModuleNotFoundError)
+    # because Python 3.10 raises ImportError for missing relative submodules.
+    # The intentional staleness ImportError is raised outside the try/except
+    # so it is never swallowed.
+    _bi = None
+    try:
+        from . import _build_info as _bi
+    except ImportError:
+        pass  # Pre-built package — no source tree, skip check
+
+    if _bi is not None:
+        from pathlib import Path
+
+        _so = Path(_rocisa.__file__)
+        # Scan rocisa sources and stinkytofu asm-IR sources (since _rocisa.so
+        # links libstinkytofu.so for the toStinkyTofuModule / emitAssembly path).
+        # Both roots are populated by CMake; an empty one signals a malformed
+        # _build_info.py. Warn (rather than scan Path("") == the CWD) and skip it,
+        # so a regression surfaces instead of silently disabling the check.
+        # Excluded from the stinkytofu scan:
+        #   - tests/         — test code is never compiled into .so
+        #   - python_module/ — only compiled into _stinkytofu.so (Python bindings)
+        #   - src/ir/logical  — logical IR is only used by _stinkytofu.so (left
+        #                       path); _rocisa.so never touches logical modules.
+        _roots = []
+        for _name, _root in (("rocisa", _bi.SOURCE_ROOT), ("stinkytofu", _bi.STINKYTOFU_SOURCE_ROOT)):
+            if _root:
+                _roots.append(Path(_root))
+            else:
+                import warnings
 
-    _so = Path(_rocisa.__file__)
-    # Scan rocisa sources and, while stinkytofu is compiled into _rocisa.so,
-    # stinkytofu sources too. STINKYTOFU_SOURCE_ROOT is removed once rocisa
-    # and stinkytofu are loaded independently.
-    # Both roots are populated by CMake; an empty one signals a malformed
-    # _build_info.py. Warn (rather than scan Path("") == the CWD) and skip it,
-    # so a regression surfaces instead of silently disabling the check.
-    _roots = []
-    for _name, _root in (("rocisa", _bi.SOURCE_ROOT), ("stinkytofu", _bi.STINKYTOFU_SOURCE_ROOT)):
-        if _root:
-            _roots.append(Path(_root))
-        else:
-            import warnings
-
-            warnings.warn(
-                f"rocisa staleness check: {_name} source root is unset in "
-                f"_build_info.py; skipping it. Rebuild with: invoke rocisa",
-                stacklevel=2,
+                warnings.warn(
+                    f"rocisa staleness check: {_name} source root is unset in "
+                    f"_build_info.py; skipping it. Rebuild with: invoke rocisa",
+                    stacklevel=2,
+                )
+        _st_root = Path(_bi.STINKYTOFU_SOURCE_ROOT) if _bi.STINKYTOFU_SOURCE_ROOT else None
+        _st_skip = {_st_root / "tests", _st_root / "src" / "ir" / "logical", _st_root / "python_module"} if _st_root else set()
+        _all_stale = _find_stale_sources(_so, _roots, _bi.BUILD_DIR)
+        _stale = [s for s in _all_stale if not any(Path(s).is_relative_to(sk) for sk in _st_skip)]
+        if _stale:
+            _preview = _stale[:3] + (["..."] if len(_stale) > 3 else [])
+            raise ImportError(
+                "rocisa C++ sources are newer than the built _rocisa.so — bindings are stale.\n"
+                f"  Modified: {', '.join(_preview)}\n"
+                "  Rebuild:  invoke rocisa"
             )
-    _stale = _find_stale_sources(_so, _roots, _bi.BUILD_DIR)
-    if _stale:
-        _preview = _stale[:3] + (["..."] if len(_stale) > 3 else [])
-        raise ImportError(
-            "rocisa C++ sources are newer than the built _rocisa.so — bindings are stale.\n"
-            f"  Modified: {', '.join(_preview)}\n"
-            "  Rebuild:  invoke rocisa"
-        )
-    del _bi, _so, _stale, _roots, _name, _root, Path
+        del _bi, _so, _stale, _all_stale, _roots, _name, _root, _st_root, _st_skip, Path
 
 
 def hasStinkyTofuBackend() -> bool: