Round-3 verification: ROCm/rocm-libraries PR #9583 @ 5bb9126a (LLVM 23.0.0git)

10_regression_diff_between_heads.sh

#!/bin/bash
# Which kernels changed between two PR heads? Instruments the suite's own _lower
# to dump IR, runs it at both heads, and diffs. Catches collateral damage from a
# "small fix" that touches many production files.
#
# Env: OLD_SRC, OLD_BUILD, NEW_SRC, NEW_BUILD, PYDEPS
: "${OLD_SRC:?}"; : "${NEW_SRC:?}"
instrument(){ python3 - "$1" "$2" <<'PY'
import pathlib, sys
d, tag = sys.argv[1], sys.argv[2]
p = pathlib.Path(d)/'tests'/'test_rocke.py'; s = p.read_text()
old = '''        b = self._builder(name)
        build(b)
        return lower_kernel_to_llvm(b.kernel, **kw)
'''
new = f'''        b = self._builder(name)
        build(b)
        ll = lower_kernel_to_llvm(b.kernel, **kw)
        import pathlib as _pl
        _o = _pl.Path("/tmp/irdump_{tag}"); _o.mkdir(parents=True, exist_ok=True)
        (_o / (name + ".ll")).write_text(ll)
        return ll
'''
assert old in s, f"anchor missing in {tag}"
p.write_text(s.replace(old, new))
PY
}
rm -rf /tmp/irdump_old /tmp/irdump_new
instrument "$OLD_SRC" old; instrument "$NEW_SRC" new
(cd "$OLD_SRC/tests" && PYTHONPATH="$PYDEPS:$OLD_SRC/python:$OLD_BUILD" python3 -m pytest test_rocke.py -q --no-header -k TestNewTargetIntrinsics >/dev/null 2>&1)
(cd "$NEW_SRC/tests" && PYTHONPATH="$PYDEPS:$NEW_SRC/python:$NEW_BUILD" python3 -m pytest test_rocke.py -q --no-header -k TestNewTargetIntrinsics >/dev/null 2>&1)
comm -12 <(ls /tmp/irdump_old|sort) <(ls /tmp/irdump_new|sort) > /tmp/common.txt
n=0; while read f; do diff -q /tmp/irdump_old/$f /tmp/irdump_new/$f >/dev/null || { echo "CHANGED: $f"; n=$((n+1)); }; done < /tmp/common.txt
echo "-> $n of $(wc -l < /tmp/common.txt) common kernels changed"

11_twospace_mangling.py

import sys, os
import os
S=os.environ["ROCKE_WORK"]
SRC=f"{S}/src"
sys.path.insert(0, f"{S}/build"); sys.path.insert(0, f"{SRC}/python")
import rocke_engine
from rocke.core.ir import IRBuilder, PtrType, I32
from rocke.core import ir_serialize
import rocke.core.lower_llvm as LL
def mk():
    b = IRBuilder("twospace"); b.kernel.attrs["max_workgroup_size"]=64
    g = b.param("g", PtrType(I32,"global"), align=4)
    c = b.param("c", PtrType(I32,"global"), align=4, addr_space="constant")
    b.s_prefetch_inst(g, b.const_i32(8))
    b.s_prefetch_inst(c, b.const_i32(8))
    return b.kernel
py  = LL._lower_kernel_to_llvm_python(mk(), arch="gfx950", llvm_flavor="llvm23")
cpp = rocke_engine.lower_serialized_ir(ir_serialize.serialize(mk()), arch="gfx950", flavor="llvm23")
open(f"{S}/twospace.py.ll","w").write(py); open(f"{S}/twospace.cpp.ll","w").write(cpp)
print("engines byte-identical:", py == cpp)
for tag, ll in (("PY ",py),("CPP",cpp)):
    ds = [l.strip() for l in ll.splitlines() if "s.prefetch.inst" in l]
    print(f"  {tag}: " + " | ".join(ds))

12_wqm_suffix_check.sh

#!/bin/bash
# s_wqm has two independently-overloaded types, so canonical mangling is .i32.i32.
# LLVM's own checked-in test uses .i32 and works only via auto-upgrade.
cd "$(mktemp -d)"
for v in "i32 i32" "i32.i32 i32" "i64 i64"; do
  set -- $v
  printf 'declare %s @llvm.amdgcn.s.wqm.%s(%s)\ndefine %s @f(%s %%m){ %%r = call %s @llvm.amdgcn.s.wqm.%s(%s %%m)\n ret %s %%r }\n' \
    "$2" "$1" "$2" "$2" "$2" "$2" "$1" "$2" "$2" > t.ll
  printf "  %-12s -> %s\n" ".$1" "$(opt -passes=verify t.ll -S -o - 2>&1 | grep -oE '@llvm\.amdgcn\.s\.wqm[a-z0-9.]*' | head -1)"
done

README.md

Round-3 verification — ROCm/rocm-libraries PR #9583 @ 5bb9126a

Verified on LLVM 23.0.0git (ROCm/llvm-project@9f1cce5a, ROCm 7.15); both engines built from source.

script	what it shows
10_regression_diff_between_heads.sh	exactly 5 of 30 kernels changed vs f19a5b09 — the five defective ones; 25 byte-identical
11_twospace_mangling.py	s_prefetch_inst on a global + a constant pointer in one module: .p1 + .p4 declares, engines byte-identical, valid IR
12_wqm_suffix_check.sh	canonical s_wqm mangling is .i32.i32; the one-suffix form works only via auto-upgrade

Also verified (not scripted here): golden IR no drift across llvm20/22/23; full TestNewTargetIntrinsics green under ROCKE_BACKEND=python|cpp|both; green with an llvm-as assertion added to every case.