Round-2 repros + measured fixes: ROCm/rocm-libraries PR #9583 @ f19a5b09 (LLVM 23.0.0git)

05_emit_both_engines.py

#!/usr/bin/env python3
"""Emit IR from BOTH rocke engines for the affected cases. Writes $ROCKE_WORK/repro2/*.ll

Env: ROCKE_WORK with subdirs  src/ (rocke/platform @ f19a5b09)  build/ (rocke_engine*.so)
"""
import os, sys
W = os.environ["ROCKE_WORK"]
sys.path.insert(0, f"{W}/build"); sys.path.insert(0, f"{W}/src/python")
import rocke_engine
from rocke.core.ir import IRBuilder, PtrType, F16, I32
from rocke.core import ir_serialize
import rocke.core.lower_llvm as LL

pyl = LL._lower_kernel_to_llvm_python          # the TRUE python lowerer
OUT = f"{W}/repro2"; os.makedirs(OUT, exist_ok=True)

def emit(name, fn, arch="gfx950"):
    def mk():
        b = IRBuilder(name); b.kernel.attrs["max_workgroup_size"] = 64; fn(b); return b.kernel
    py  = pyl(mk(), arch=arch, llvm_flavor="llvm23")
    cpp = rocke_engine.lower_serialized_ir(ir_serialize.serialize(mk()), arch=arch, flavor="llvm23")
    open(f"{OUT}/{name}.py.ll", "w").write(py)
    open(f"{OUT}/{name}.cpp.ll", "w").write(cpp)
    print(f"  {name:26s} engines byte-identical: {py == cpp}")

def av_load(b):  b.av_load_b128(b.param("p", PtrType(I32, "global"), align=16))
def av_store(b):
    p = b.param("p", PtrType(I32, "global"), align=16); b.av_store_b128(p, b.av_load_b128(p))
def prefetch(b): b.s_prefetch_inst(b.param("code", PtrType(I32, "global"), align=4), b.const_i32(64))
def bufasync(b):
    X = b.param("X", PtrType(F16, "global")); N = b.param("N", I32)
    lds = b.smem_alloc(F16, [64, 8], name_hint="stage")
    b.buffer_load_lds_async(b.buffer_rsrc(X, N), b.smem_addr_of(lds),
                            b.const_i32(0), b.const_i32(0), dwords=4)
def glasync(b):
    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)
    b.global_load_async_to_lds(X, z, c, [z, z], width_bytes=16)
    b.smem_load_v4_f16(c, z, z)

for n, f in [("av_load_b128", av_load), ("av_store_b128", av_store),
             ("s_prefetch_inst", prefetch), ("buffer_load_lds_async", bufasync),
             ("global_load_async_to_lds", glasync)]:
    try: emit(n, f)
    except Exception as e: print(f"  {n:26s} EXC {type(e).__name__}: {str(e)[:80]}")

06_validate_on_llvm23.sh

#!/bin/bash
# Emit from both engines, then validate every module on LLVM 23.
: "${ROCKE_WORK:?set ROCKE_WORK (needs src/ and build/ subdirs)}"
PYDEPS="${PYDEPS:-$ROCKE_WORK/pydeps}"
echo "toolchain: $(clang --version 2>/dev/null | head -1)"
echo; echo "########## A. generate IR from both engines ##########"
PYTHONPATH="$PYDEPS" python3 "$(dirname "$0")/05_emit_both_engines.py" || exit 1
echo; echo "########## B. validate each on LLVM 23 ##########"
for f in "$ROCKE_WORK"/repro2/*.ll; do
  out=$(opt -passes=verify "$f" -S -o /dev/null 2>&1)
  if [ $? -eq 0 ]; then printf "  OK       %s\n" "$(basename "$f")"
  else printf "  REJECTED %-32s %s\n" "$(basename "$f")" "$(echo "$out" | grep -m1 -o 'error:.*' | cut -c1-70)"; fi
done

07_tests_pass_on_invalid_ir.sh

#!/bin/bash
# The PR's own tests, run against the exact code that emits the rejected IR.
: "${ROCKE_WORK:?set ROCKE_WORK (needs src/ and build/ subdirs)}"
PYDEPS="${PYDEPS:-$ROCKE_WORK/pydeps}"
export PYTHONPATH="$PYDEPS"
python3 -c "import pytest" 2>/dev/null || pip3 install --target="$PYDEPS" --no-cache-dir pytest >/dev/null 2>&1
echo "pytest $(python3 -c 'import pytest;print(pytest.__version__)')"
cd "$ROCKE_WORK/src/tests" || exit 1
PYTHONPATH="$PYDEPS:$ROCKE_WORK/src/python:$ROCKE_WORK/build" python3 -m pytest test_rocke.py \
  -q --no-header -p no:cacheprovider \
  -k "av_load_b128_emits or av_store_b128_emits or s_prefetch_inst_emits or buffer_load_lds_async_converts" \
  2>&1 | tail -5

08_suggested_fix_validate_ir.patch

--- a/tests/test_rocke.py	2026-07-28 21:57:29.000000000 -0400
+++ b/tests/test_rocke.py	2026-07-28 23:06:15.000000000 -0400
@@ -2731,13 +2731,24 @@
         b.kernel.attrs["max_workgroup_size"] = 64
         return b
 
+    def _assert_ir_assembles(self, name, ll):
+        import os, subprocess, shutil
+        tool = shutil.which("llvm-as") or "/opt/rocm/llvm/bin/llvm-as"
+        if not os.path.exists(tool):
+            return
+        r = subprocess.run([tool, "-o", os.devnull], input=ll, text=True, capture_output=True)
+        if r.returncode != 0:
+            self.fail(f"{name}: emitted IR rejected by llvm-as:\n{r.stderr.strip()}")
+
     def _lower(self, name, build, **kw):
         """Build a single-intrinsic kernel and return its lowered LLVM IR."""
         from rocke.core.lower_llvm import lower_kernel_to_llvm
 
         b = self._builder(name)
         build(b)
-        return lower_kernel_to_llvm(b.kernel, **kw)
+        ll = lower_kernel_to_llvm(b.kernel, **kw)
+        self._assert_ir_assembles(name, ll)
+        return ll
 
     # ---- ds_swizzle (raw offset + XOR-butterfly encoding) ----
     def test_ds_swizzle_passes_raw_offset_immediate(self):

09_measure_fix_effect.sh

#!/bin/bash
# Measures what each suggested change catches. Needs $ROCKE_WORK/src and $ROCKE_WORK/build.
: "${ROCKE_WORK:?set ROCKE_WORK}"; PYDEPS="${PYDEPS:-$ROCKE_WORK/pydeps}"
cd "$ROCKE_WORK/src/tests" || exit 1
PP="$PYDEPS:$ROCKE_WORK/src/python:$ROCKE_WORK/build"
for be in "" python cpp both; do
  printf "ROCKE_BACKEND=%-8s " "${be:-unset}"
  env ${be:+ROCKE_BACKEND=$be} PYTHONPATH="$PP" python3 -m pytest test_rocke.py \
    -q --no-header -p no:cacheprovider -k TestNewTargetIntrinsics 2>&1 | tail -1
done
echo
echo "# then apply 08_suggested_fix_validate_ir.patch and re-run: expect 5 failed / 27 passed"

README.md

Round-2 repros — ROCm/rocm-libraries PR #9583 @ f19a5b09

Verified on LLVM 23.0.0git (ROCm/llvm-project@9f1cce5a); engine built from PR head.

export ROCKE_WORK=/some/scratch     # needs:
#   $ROCKE_WORK/src    = <checkout>/dnn-providers/hip-kernel-provider/rocke/platform @ f19a5b09
#   $ROCKE_WORK/build  = dir containing rocke_engine*.so   (build: round-1 gist, 01_build_engine.sh)
bash 06_validate_on_llvm23.sh          # emit from both engines, validate each
bash 07_tests_pass_on_invalid_ir.sh    # the PR's own tests, same code
bash 09_measure_fix_effect.sh          # what each suggested change catches

What the emitted IR does

case	engines agree?	LLVM 23
av_load_b128	identical	REJECTED '%p' ... 'ptr addrspace(1)' but expected 'ptr'
av_store_b128	identical	REJECTED (same)
s_prefetch_inst	identical	REJECTED (same)
buffer_load_lds_async	identical	REJECTED '%lds_addr3' ... 'i64' but expected 'ptr addrspace(3)'
global_load_async_to_lds	DIFFER	C++ REJECTED use of undefined value '@stageB3...'; Python OK

…while the PR's own tests report 4 passed.

Two changes, measured

change	cost	catches
run the suite once with ROCKE_BACKEND=both	no code change	the C++/Python divergence — 1 failed / 31 passed, flags exactly test_global_load_async_to_lds_supports_b8_width
08_suggested_fix_validate_ir.patch — llvm-as check in the shared _lower helper	~10 lines, one function	all 5 defective cases; 27/32 still pass, no false positives

Baseline for comparison: ROCKE_BACKEND unset / python / cpp all report 32 passed.

The patch skips silently when llvm-as is absent, so it is safe in environments without a toolchain — though note it does introduce subprocess, which test_rocke.py:15 currently disclaims ("no subprocess, no GPU"); gating it behind an env var would keep that contract.