Repros for review of ROCm/rocm-libraries PR #9583 (rocke LLVM-23 intrinsics) — verified on LLVM 23.0.0git

01_build_engine.sh

#!/bin/bash
# Build the rocke C++ engine + pybind extension from PR #9583 head.
# Env: ROCKE_SRC = <checkout>/dnn-providers/hip-kernel-provider/rocke/platform
#      OUT       = build output dir
set -eo pipefail
ROCKE_SRC="${ROCKE_SRC:?set ROCKE_SRC to .../rocke/platform}"
OUT="${OUT:-/tmp/rocke_build}"
PYDEPS="${PYDEPS:-/tmp/rocke_pydeps}"

export PYTHONPATH="$PYDEPS:$PYTHONPATH"
python3 -c "import pybind11" 2>/dev/null || pip3 install --target="$PYDEPS" --no-cache-dir "pybind11>=3.0"
PB_DIR=$(python3 -c "import pybind11; print(pybind11.get_cmake_dir())")

cmake -S "$ROCKE_SRC" -B "$OUT/engine" -G Ninja -DCMAKE_BUILD_TYPE=Release
cmake --build "$OUT/engine" --target rocke_core -j"$(nproc)"

cmake -S "$ROCKE_SRC/cpp/bindings" -B "$OUT/build" -G Ninja \
  -DCMAKE_BUILD_TYPE=Release \
  -DROCKE_ENGINE_ARCHIVE="$OUT/engine/librocke_core.a" \
  -Dpybind11_DIR="$PB_DIR" \
  -DPYTHON_EXECUTABLE="$(command -v python3)"
cmake --build "$OUT/build" -j"$(nproc)"
find "$OUT/build" -name "rocke_engine*.so"

02_finding1_lds_base.py

#!/usr/bin/env python3
"""Finding 1: op_tile_global_load_async_to_lds uses the raw smem gname as the
GEP base instead of the pool-relative base pointer.

Calls the C++ engine DIRECTLY (bypassing backend.py's flavor gate, see finding 2).
Env: ROCKE_SRC, ROCKE_BUILD (dir containing rocke_engine*.so)
"""
import os, sys
SRC   = os.environ["ROCKE_SRC"]
BUILD = os.environ["ROCKE_BUILD"]
sys.path.insert(0, BUILD); sys.path.insert(0, f"{SRC}/python")

import rocke_engine
from rocke.core.ir import IRBuilder, PtrType, F16
from rocke.core import ir_serialize
from rocke.core.lower_llvm import lower_kernel_to_llvm as py_lower

def build():
    b = IRBuilder("f2_repro"); b.kernel.attrs["max_workgroup_size"] = 64
    X = b.param("X", PtrType(F16, "global")); z = b.const_i32(0)
    a = b.smem_alloc(F16, [64, 8], name_hint="stageA")
    c = b.smem_alloc(F16, [64, 8], name_hint="stageB")
    b.smem_load_v4_f16(a, z, z)                                  # keep A live -> offset 0
    b.global_load_async_to_lds(X, z, c, [z, z], width_bytes=16)  # async copy into B
    b.smem_load_v4_f16(c, z, z)                                  # keep B live -> offset 1024
    return b.kernel

cpp_ll = rocke_engine.lower_serialized_ir(ir_serialize.serialize(build()),
                                          arch="gfx950", flavor="llvm23")
py_ll  = py_lower(build(), arch="gfx950", llvm_flavor="llvm23")
open("f2_ENGINE_cpp.ll", "w").write(cpp_ll)
open("f2_ENGINE_py.ll",  "w").write(py_ll)

print("byte-identical?", cpp_ll == py_ll)
for tag, ll in (("CPP", cpp_ll), ("PY ", py_ll)):
    print(f"--- {tag} ---")
    for l in ll.splitlines():
        if ("async_dst" in l and "getelementptr" in l) or (l.startswith("@") and "addrspace(3)" in l):
            print("   ", l.strip())
print("\nNow: opt -passes=verify f2_ENGINE_cpp.ll   # expect: use of undefined value")

03_finding2_backend_flavor_gate.py

#!/usr/bin/env python3
"""Finding 2: python/rocke/core/backend.py still gates the cpp path on
(LLVM20, LLVM22), so on an llvm23 host the default cpp backend silently
falls back to the Python lowerer.

Env: ROCKE_SRC, ROCKE_BUILD. Run on a ROCm >= 7.13 host (flavor -> llvm23).
"""
import os, sys
SRC   = os.environ["ROCKE_SRC"]
BUILD = os.environ["ROCKE_BUILD"]
sys.path.insert(0, BUILD); sys.path.insert(0, f"{SRC}/python")

from rocke.core.lower_llvm import _resolve_llvm_flavor, _detect_llvm_flavor
from rocke.core import backend as bk
print("ROCm-detected flavor :", _detect_llvm_flavor())
print("resolved flavor      :", _resolve_llvm_flavor())
print("default backend      :", bk._DEFAULT_BACKEND)

from rocke.core.ir import IRBuilder, PtrType, F16
from rocke.core.lower_llvm import lower_kernel_to_llvm
b = IRBuilder("k"); b.kernel.attrs["max_workgroup_size"] = 64
X = b.param("X", PtrType(F16, "global")); z = b.const_i32(0)
s = b.smem_alloc(F16, [64, 8], name_hint="s"); b.smem_load_v4_f16(s, z, z)

print("\n--- lowering with DEFAULT backend, no overrides ---")
ll = lower_kernel_to_llvm(b.kernel, arch="gfx950")
print("lowered OK, len", len(ll), "(note the fallback WARNING above)")

04_finding3_prefetch_addrspace.sh

#!/bin/bash
# Finding 3: s_prefetch_inst is emitted with a hardcoded addrspace(0) `ptr`.
# Needs an LLVM 23 `opt` on PATH.
cd "$(mktemp -d)"
chk(){ printf "%-46s" "$1"
  out=$(opt -passes=verify "$2" -S -o /dev/null 2>&1)
  [ $? -eq 0 ] && echo "ACCEPTED" || echo "REJECTED: $(echo "$out"|grep -m1 -i error|cut -c1-88)"; }

cat > as0.ll <<'X'
declare void @llvm.amdgcn.s.prefetch.inst(ptr, i32)
define void @f(ptr %p, i32 %n) { call void @llvm.amdgcn.s.prefetch.inst(ptr %p, i32 %n)
  ret void }
X
cat > as4.ll <<'X'
declare void @llvm.amdgcn.s.prefetch.inst(ptr, i32)
define void @f(ptr addrspace(4) %p, i32 %n) { call void @llvm.amdgcn.s.prefetch.inst(ptr %p, i32 %n)
  ret void }
X
cat > as1.ll <<'X'
declare void @llvm.amdgcn.s.prefetch.inst(ptr, i32)
define void @f(ptr addrspace(1) %p, i32 %n) { call void @llvm.amdgcn.s.prefetch.inst(ptr %p, i32 %n)
  ret void }
X
chk "operand addrspace(0)  (PR emits 'ptr')" as0.ll
chk "operand addrspace(4)  (instruction mem)" as4.ll
chk "operand addrspace(1)  (global)" as1.ll
echo
echo "NB: the missing .pN suffix is NOT a defect -- LLVM auto-upgrades it:"
opt -passes=verify as0.ll -S -o - 2>/dev/null | grep -E "^\s*call"

README.md

Repros for review of ROCm/rocm-libraries PR #9583

Environment used: container rocm-devtools-therock-7.15 (LLVM 23.0.0git, Python 3.12), engine built from PR head 56006f38.

export ROCKE_SRC=<checkout>/dnn-providers/hip-kernel-provider/rocke/platform
export OUT=/tmp/rocke_build
./01_build_engine.sh
export ROCKE_BUILD=$OUT/build
python3 02_finding1_lds_base.py && opt -passes=verify f2_ENGINE_cpp.ll
python3 03_finding2_backend_flavor_gate.py
bash    04_finding3_prefetch_addrspace.sh

file	finding
02_finding1_lds_base.py	global_load_async_to_lds C++ uses raw smem gname as GEP base -> undefined global + byte-identity break
03_finding2_backend_flavor_gate.py	backend.py rejects llvm23 -> default cpp backend silently falls back to Python
04_finding3_prefetch_addrspace.sh	s_prefetch_inst hardcodes addrspace(0)