amd isel renderer

* renderer/amd: add cdna emulator

* fixes

* no predecode

* no early

* REMU_PATH

* delete that

* round

* Fix cache invalidation check in _compile_smem

.github/workflows/test.yml

@@ -708,6 +708,8 @@ jobs:
         run: SKIP_SLOW_TEST=1 AMD_LLVM=0 pytest -n=auto test/backend/test_ops.py -k "test_sparse_categorical_crossentropy or test_tril or test_nonzero or test_softmax_argmax" --durations 20
       - name: Run RDNA4 emulator tests
         run: MOCKGPU_ARCH=rdna4 python -m pytest test/test_tiny.py -v --durations 20
+      - name: Run CDNA4 emulator tests
+        run: AMD_LLVM=1 MOCKGPU_ARCH=cdna4 python -m pytest test/test_tiny.py -v --durations 20
 
   testnvidia:
     strategy:

extra/mmapeak/mmapeak.py

@@ -4,15 +4,16 @@ import os
 os.environ["AMD_AQL"] = "1"
 
 from tinygrad import Tensor, Device
+from tinygrad.helpers import getenv
 from tinygrad.uop.ops import UOp, Ops, KernelInfo
 from tinygrad.renderer import Estimates
 from tinygrad.renderer.amd.dsl import Reg, Inst, s, v
 
 NUM_WORKGROUPS = 96
 WAVE_SIZE = 32
-NUM_WAVES = 2
+NUM_WAVES = 4
 FLOPS_PER_MATMUL = 16*16*16*2
-INTERNAL_LOOP = 1_000_00
+INTERNAL_LOOP = getenv("LOOP", 10_000)
 INSTRUCTIONS_PER_LOOP = 200
 
 def repeat(insts:list[Inst], n:int, counter_sreg:Reg) -> list[Inst]:
@@ -22,15 +23,6 @@ def repeat(insts:list[Inst], n:int, counter_sreg:Reg) -> list[Inst]:
   branch_inst = s_cbranch_scc1(simm16=-((loop_sz // 4) + 1) & 0xFFFF)
   return [s_mov_b32(counter_sreg, n)] + insts + [sub_inst, cmp_inst, branch_inst, s_endpgm()]
 
-def make_kernel(insts:list[Inst]):
-  def fxn(A:UOp) -> UOp:
-    threads = UOp.special(WAVE_SIZE * NUM_WAVES, "lidx0")
-    gidx = UOp.special(NUM_WORKGROUPS, "gidx0")
-    FLOPs = FLOPS_PER_MATMUL * NUM_WAVES * NUM_WORKGROUPS * INTERNAL_LOOP * INSTRUCTIONS_PER_LOOP
-    sink = UOp.sink(A.base, threads, gidx, arg=KernelInfo("mmapeak", estimates=Estimates(ops=FLOPs, mem=0)))
-    return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg="AMD"), UOp(Ops.LINEAR, src=tuple([UOp(Ops.INS, arg=x) for x in insts]))))
-  return fxn
-
 def launchBenchmark(instruction, vgprIndices, dense=True, accum=False, **kwargs):
   if accum:
     inst = instruction(v[0:vgprIndices[0]], v[vgprIndices[1]:vgprIndices[2]], v[vgprIndices[1]:vgprIndices[2]], 1, acc_cd=1, **kwargs)
@@ -39,7 +31,12 @@ def launchBenchmark(instruction, vgprIndices, dense=True, accum=False, **kwargs)
   else:
     inst = instruction(v[0:vgprIndices[0]], v[vgprIndices[1]:vgprIndices[2]], v[vgprIndices[3]:vgprIndices[4]], v[vgprIndices[5]])
   insts = repeat([inst for _ in range(INSTRUCTIONS_PER_LOOP)], n=INTERNAL_LOOP, counter_sreg=s[1])
-  fxn = make_kernel(insts)
+  def fxn(A:UOp) -> UOp:
+    threads = UOp.special(WAVE_SIZE * NUM_WAVES, "lidx0")
+    gidx = UOp.special(NUM_WORKGROUPS, "gidx0")
+    FLOPs = FLOPS_PER_MATMUL * NUM_WAVES * NUM_WORKGROUPS * INTERNAL_LOOP * INSTRUCTIONS_PER_LOOP
+    sink = UOp.sink(A.base, threads, gidx, arg=KernelInfo(inst.op.name.lower(), estimates=Estimates(ops=FLOPs, mem=0)))
+    return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg="AMD"), UOp(Ops.LINEAR, src=tuple([UOp(Ops.INS, arg=x) for x in insts]))))
   dummy = Tensor.zeros(1).contiguous().realize()
   out = Tensor.custom_kernel(dummy, fxn=fxn)[0]
   ei = out.schedule()[-1].lower()

test/amd/bench_emu.py

@@ -1,267 +0,0 @@
-#!/usr/bin/env python3
-"""Benchmark comparing Python vs Rust RDNA3 emulators on real tinygrad kernels."""
-import ctypes, time, os
-from pathlib import Path
-
-from tinygrad.renderer.amd.emu import run_asm as python_run_asm, decode_program
-from tinygrad.renderer.amd import decode_inst
-from tinygrad.runtime.autogen.amd.rdna3.ins import SOPP, SOPPOp
-
-import tinygrad
-EXTRA_DIR = Path(tinygrad.__file__).parent.parent / "extra"
-REMU_PATH = EXTRA_DIR / "remu/target/release/libremu.so"
-if not REMU_PATH.exists():
-  REMU_PATH = EXTRA_DIR / "remu/target/release/libremu.dylib"
-
-def get_rust_remu():
-  """Load the Rust libremu shared library."""
-  if not REMU_PATH.exists(): return None
-  remu = ctypes.CDLL(str(REMU_PATH))
-  remu.run_asm.restype = ctypes.c_int32
-  remu.run_asm.argtypes = [ctypes.c_void_p, ctypes.c_uint32, ctypes.c_uint32, ctypes.c_uint32, ctypes.c_uint32,
-                           ctypes.c_uint32, ctypes.c_uint32, ctypes.c_uint32, ctypes.c_void_p]
-  return remu
-
-def count_instructions(kernel: bytes) -> int:
-  """Count instructions in a kernel."""
-  return len(decode_program(kernel))
-
-def setup_buffers(buf_sizes: list[int], init_data: dict[int, bytes] | None = None):
-  """Allocate buffers and return args pointer + valid ranges."""
-  if init_data is None: init_data = {}
-  buffers = []
-  for i, size in enumerate(buf_sizes):
-    padded = ((size + 15) // 16) * 16 + 16
-    data = init_data.get(i, b'\x00' * padded)
-    data_list = list(data) + [0] * (padded - len(data))
-    buf = (ctypes.c_uint8 * padded)(*data_list[:padded])
-    buffers.append(buf)
-  args = (ctypes.c_uint64 * len(buffers))(*[ctypes.addressof(b) for b in buffers])
-  args_ptr = ctypes.addressof(args)
-  ranges = {(ctypes.addressof(b), len(b)) for b in buffers}
-  ranges.add((args_ptr, ctypes.sizeof(args)))
-  return buffers, args, args_ptr, ranges
-
-def benchmark_emulator(name: str, run_fn, kernel: bytes, global_size, local_size, args_ptr, rsrc2: int, iterations: int = 5):
-  """Benchmark an emulator and return average time."""
-  gx, gy, gz = global_size
-  lx, ly, lz = local_size
-  kernel_buf = (ctypes.c_char * len(kernel)).from_buffer_copy(kernel)
-  lib_ptr = ctypes.addressof(kernel_buf)
-
-  # Warmup
-  run_fn(lib_ptr, len(kernel), gx, gy, gz, lx, ly, lz, args_ptr, rsrc2)
-
-  # Timed runs
-  times = []
-  for _ in range(iterations):
-    start = time.perf_counter()
-    result = run_fn(lib_ptr, len(kernel), gx, gy, gz, lx, ly, lz, args_ptr, rsrc2)
-    end = time.perf_counter()
-    if result != 0:
-      print(f"  {name} returned error: {result}")
-      return None
-    times.append(end - start)
-
-  return sum(times) / len(times)
-
-def profile_instructions(kernel: bytes):
-  """Profile individual instruction compile times."""
-  from tinygrad.renderer.amd.emu import _get_runner, _canonical_runner_cache
-  from tinygrad.helpers import Context
-  _get_runner.cache_clear()
-  _canonical_runner_cache.clear()
-
-  results = []
-  i = 0
-  while i < len(kernel):
-    inst = decode_inst(kernel[i:])
-    if isinstance(inst, SOPP) and inst.op == SOPPOp.S_CODE_END: break
-    inst_bytes = bytes(kernel[i:i + inst.size() + 4])
-    try: inst_str = repr(inst)
-    except Exception: inst_str = f"<{type(inst).__name__}>"
-
-    # Time the full compile (sink + render + compile)
-    start = time.perf_counter()
-    with Context(CCACHE=0):
-      runner, is_new = _get_runner(inst_bytes)
-    compile_time = time.perf_counter() - start
-
-    results.append({
-      'inst_str': inst_str + ('' if is_new else ' [CACHED]'),
-      'compile_ms': compile_time * 1000 if is_new else 0,
-    })
-    i += inst.size()
-
-  return sorted(results, key=lambda x: x['compile_ms'], reverse=True)
-
-def benchmark_python_split(kernel: bytes, global_size, local_size, args_ptr, rsrc2: int, iterations: int = 5):
-  """Benchmark Python emulator with compile and execution times."""
-  from tinygrad.renderer.amd.emu import _get_runner, _canonical_runner_cache
-  from tinygrad.helpers import Context
-  _get_runner.cache_clear()
-  _canonical_runner_cache.clear()
-  decode_program.cache_clear()
-
-  # Measure compile time (decode_program builds sinks, renders, and compiles)
-  compile_start = time.perf_counter()
-  with Context(CCACHE=0):
-    program = decode_program(kernel)
-  compile_time = time.perf_counter() - compile_start
-  n_compiled = len(_canonical_runner_cache)
-
-  # Execution time
-  exec_time = benchmark_emulator("Python", python_run_asm, kernel, global_size, local_size, args_ptr, rsrc2, iterations)
-  return compile_time, exec_time, len(program), n_compiled
-
-def get_tinygrad_kernel(op_name: str) -> tuple[bytes, tuple, tuple, list[int], dict[int, bytes], int] | None:
-  """Get a real tinygrad kernel by operation name. Returns (code, global_size, local_size, buf_sizes, buf_data, rsrc2)."""
-  try:
-    from tinygrad import Tensor
-    from tinygrad.runtime.support.elf import elf_loader
-    from tinygrad.runtime.autogen import hsa
-    import numpy as np
-    np.random.seed(42)
-
-    ops = {
-      "add": lambda: Tensor.empty(1024) + Tensor.empty(1024),
-      "mul": lambda: Tensor.empty(1024) * Tensor.empty(1024),
-      "matmul_small": lambda: Tensor.empty(16, 16) @ Tensor.empty(16, 16),
-      "matmul_medium": lambda: Tensor.empty(64, 64) @ Tensor.empty(64, 64),
-      "reduce_sum": lambda: Tensor.empty(4096).sum(),
-      "reduce_max": lambda: Tensor.empty(4096).max(),
-      "softmax": lambda: Tensor.empty(256).softmax(),
-      "layernorm": lambda: Tensor.empty(32, 64).layernorm(),
-      "conv2d": lambda: Tensor.empty(1, 4, 16, 16).conv2d(Tensor.empty(4, 4, 3, 3)),
-      "gelu": lambda: Tensor.empty(1024).gelu(),
-      "exp": lambda: Tensor.empty(1024).exp(),
-      "sin": lambda: Tensor.empty(1024).sin(),
-    }
-
-    if op_name not in ops: return None
-    out = ops[op_name]()
-    sched = out.schedule()
-
-    for ei in sched:
-      lowered = ei.lower()
-      if ei.ast.op.name == 'SINK' and lowered.prg and lowered.prg.p.lib:
-        lib = bytes(lowered.prg.p.lib)
-        image = memoryview(bytearray(lib))
-        _, sections, _ = elf_loader(lib)
-        rodata_entry = next((sh.header.sh_addr for sh in sections if sh.name == ".rodata"), -1)
-        for sec in sections:
-          if sec.name == '.text':
-            buf_sizes = [b.nbytes for b in lowered.bufs]
-            # Get initial data from numpy arrays if available
-            buf_data = {}
-            for i, buf in enumerate(lowered.bufs):
-              if hasattr(buf, 'base') and buf.base is not None and hasattr(buf.base, '_buf'):
-                try: buf_data[i] = bytes(buf.base._buf)
-                except Exception: pass
-            # Extract rsrc2 from ELF (same as ops_amd.py)
-            group_segment_size = image[rodata_entry:rodata_entry+4].cast("I")[0]
-            lds_size = ((group_segment_size + 511) // 512) & 0x1FF
-            code = hsa.amd_kernel_code_t.from_buffer_copy(bytes(image[rodata_entry:rodata_entry+256]) + b'\x00'*256)
-            rsrc2 = code.compute_pgm_rsrc2 | (lds_size << 15)
-            return (bytes(sec.content), tuple(lowered.prg.p.global_size), tuple(lowered.prg.p.local_size), buf_sizes, buf_data, rsrc2)
-    return None
-  except Exception as e:
-    print(f"  Error getting kernel: {e}")
-    return None
-
-TINYGRAD_TESTS = ["add", "mul", "reduce_sum", "softmax", "exp", "sin", "gelu", "matmul_small"]
-
-def main():
-  import argparse
-  parser = argparse.ArgumentParser(description="Benchmark RDNA3 emulators")
-  parser.add_argument("--iterations", type=int, default=3, help="Number of iterations per benchmark")
-  parser.add_argument("--profile", type=str, default=None, help="Profile instructions for a specific kernel (e.g. 'sin')")
-  parser.add_argument("--top", type=int, default=20, help="Number of top instructions to show in profile")
-  args = parser.parse_args()
-
-  # Profile mode: show individual instruction timing
-  if args.profile:
-    kernel_info = get_tinygrad_kernel(args.profile)
-    if kernel_info is None:
-      print(f"Failed to get kernel for '{args.profile}'")
-      return
-    kernel = kernel_info[0]
-    print(f"Profiling instructions for '{args.profile}' kernel...")
-    print("=" * 110)
-    results = profile_instructions(kernel)
-    print(f"{'Instruction':<90} {'Compile(ms)':>12}")
-    print("-" * 110)
-    for r in results[:args.top]:
-      inst = r['inst_str'][:87] + "..." if len(r['inst_str']) > 90 else r['inst_str']
-      print(f"{inst:<90} {r['compile_ms']:>12.3f}")
-    print("-" * 110)
-    total = sum(r['compile_ms'] for r in results)
-    print(f"{'TOTAL':<90} {total:>12.3f}")
-    return
-
-  rust_remu = get_rust_remu()
-  if rust_remu is None:
-    print("Rust libremu not found. Build with: cargo build --release --manifest-path extra/remu/Cargo.toml")
-    print("Running Python-only benchmarks...\n")
-
-  print("=" * 90)
-  print("RDNA3 Emulator Benchmark: Python vs Rust")
-  print("=" * 90)
-
-  results = []
-
-  print("\n[TINYGRAD KERNELS]")
-  print("-" * 90)
-
-  for op_name in TINYGRAD_TESTS:
-    print(f"\n{op_name}:", end=" ", flush=True)
-    kernel_info = get_tinygrad_kernel(op_name)
-    if kernel_info is None:
-      print("failed to compile")
-      continue
-
-    kernel, global_size, local_size, buf_sizes, buf_data, rsrc2 = kernel_info
-    buffers, args_arr, args_ptr, ranges = setup_buffers(buf_sizes, buf_data)
-
-    # Benchmark Python emulator (must be first to measure compile time before cache is populated)
-    py_compile, py_exec, n_insts, n_compiled = benchmark_python_split(kernel, global_size, local_size, args_ptr, rsrc2, args.iterations)
-
-    n_workgroups = global_size[0] * global_size[1] * global_size[2]
-    n_threads = local_size[0] * local_size[1] * local_size[2]
-    total_work = n_insts * n_workgroups * n_threads
-
-    print(f"{n_insts} insts ({n_compiled} unique) × {n_workgroups} WGs × {n_threads} threads = {total_work:,} ops")
-    rust_time = benchmark_emulator("Rust", rust_remu.run_asm, kernel, global_size, local_size,
-                                   args_ptr, rsrc2, args.iterations) if rust_remu else None
-
-    if py_compile is not None:
-      py_exec_rate = total_work / py_exec / 1e6
-      print(f"  Compile:        {py_compile*1000:8.3f} ms  ({n_compiled} unique)")
-      print(f"  Exec:           {py_exec*1000:8.3f} ms  ({py_exec_rate:7.2f} M ops/s)")
-    if rust_time:
-      rust_rate = total_work / rust_time / 1e6
-      speedup = py_exec / rust_time if py_exec else 0
-      print(f"  Rust:           {rust_time*1000:8.3f} ms  ({rust_rate:7.2f} M ops/s)  [{speedup:.1f}x faster]")
-
-    results.append((op_name, n_insts, n_compiled, n_workgroups, py_compile, py_exec, rust_time))
-
-  # Summary table
-  print("\n" + "=" * 110)
-  print("SUMMARY")
-  print("=" * 110)
-  print(f"{'Name':<16} {'Insts':<6} {'Unique':<6} {'WGs':<5} {'Compile (ms)':<14} {'Exec (ms)':<12} {'Rust (ms)':<12} {'Speedup':<10}")
-  print("-" * 110)
-
-  for name, n_insts, n_compiled, n_wgs, py_compile, py_exec, rust_time in results:
-    compile_ms = f"{py_compile*1000:.3f}" if py_compile else "error"
-    exec_ms = f"{py_exec*1000:.3f}" if py_exec else "error"
-    if rust_time:
-      rust_ms = f"{rust_time*1000:.3f}"
-      speedup = f"{py_exec/rust_time:.1f}x" if py_exec else "N/A"
-    else:
-      rust_ms, speedup = "N/A", "N/A"
-    print(f"{name:<16} {n_insts:<6} {n_compiled:<6} {n_wgs:<5} {compile_ms:<14} {exec_ms:<12} {rust_ms:<12} {speedup:<10}")
-
-if __name__ == "__main__":
-  os.environ["AMD"] = "1"
-  main()

test/amd/test_compare_emulators.py

@@ -1,12 +1,15 @@
 # Test to compare Python and Rust RDNA3 emulators by running real tinygrad kernels
 import unittest, ctypes
 from dataclasses import dataclass
+from pathlib import Path
 from tinygrad import Device
 
-from tinygrad.renderer.amd.emu import WaveState, decode_program, WAVE_SIZE, VCC_LO, EXEC_LO, SCC
+from tinygrad.renderer.amd.emu import WaveState, _decode_at, WAVE_SIZE, VCC_LO, EXEC_LO, SCC
 from tinygrad.renderer.amd import decode_inst
 from test.amd.helpers import KernelInfo
-from test.amd.bench_emu import REMU_PATH
+import tinygrad
+REMU_PATH = Path(tinygrad.__file__).parent.parent / "extra/remu/target/release/libremu.so"
+if not REMU_PATH.exists(): REMU_PATH = Path(tinygrad.__file__).parent.parent / "extra/remu/target/release/libremu.dylib"
 
 def set_valid_mem_ranges(ranges): pass  # emu2 doesn't need this
 
@@ -89,7 +92,7 @@ class RustEmulator:
 class PythonEmulator:
   def __init__(self):
     self.state: WaveState | None = None
-    self.program: dict | None = None
+    self.program: dict[int, tuple] = {}  # lazily populated: pc -> (name, fxn, globals)
     self.vmem_buf = None
     self.lds_buf = None
     self.kernel_buf = None  # Keep kernel bytes alive
@@ -99,27 +102,29 @@ class PythonEmulator:
     import ctypes
     from tinygrad.device import Buffer, BufferSpec
     from tinygrad.dtype import dtypes
-    # Store kernel in a ctypes buffer so generic instructions can read from vmem at actual PC address
+    # Store kernel in a ctypes buffer so _decode_at can read from memory at actual PC address
     self.kernel_buf = (ctypes.c_char * len(kernel)).from_buffer_copy(kernel)
     self.lib_addr = ctypes.addressof(self.kernel_buf)
-    # Remap program dict to use actual addresses (like run_asm does)
-    program_raw = decode_program(kernel)
-    self.program = {self.lib_addr + offset: val for offset, val in program_raw.items()}
+    self.program = {}
     self.state = WaveState(n_lanes)
     self.state.pc = self.lib_addr  # Set PC to code base address
     self.vmem_buf = Buffer('CPU', 1 << 40, dtypes.uint32, options=BufferSpec(external_ptr=0)).ensure_allocated()
     self.lds_buf = Buffer('CPU', 65536 // 4, dtypes.uint32).ensure_allocated()
 
+  def _ensure_decoded(self, pc: int):
+    if pc not in self.program:
+      runner = _decode_at(pc, "rdna3")
+      self.program[pc] = (runner.p.function_name, runner._prg.fxn, runner.p.globals)
+
   def step(self) -> int:
     import ctypes
-    assert self.program is not None and self.state is not None
+    assert self.state is not None
     pc = self.state.pc
-    if pc == 0xFFFFFFFFFFFFFFFF or pc not in self.program: return -1
-    name, fxn, globals_list, _runner = self.program[pc]
-    if fxn is None: return 1  # unsupported instruction
+    if pc == 0xFFFFFFFFFFFFFFFF: return -1
+    self._ensure_decoded(pc)
+    name, fxn, globals_list = self.program[pc]
     buf_addrs = {0: self.state.sgpr_buf._buf.va_addr, 1: self.state.vgpr_buf._buf.va_addr,  # type: ignore[union-attr]
                  2: self.vmem_buf._buf.va_addr, 3: self.lds_buf._buf.va_addr}  # type: ignore[union-attr]
-    # Direct ctypes call - bypasses HCQ overhead
     fxn(*[ctypes.c_uint64(buf_addrs[g]) for g in globals_list], ctypes.c_int32(0))
     return -1 if self.state.pc == 0xFFFFFFFFFFFFFFFF else 0
 
@@ -140,7 +145,7 @@ class PythonEmulator:
                          exec_mask=sgpr[EXEC_LO.offset], sgpr=sgpr, vgpr=vgpr)
 
 def run_single_kernel(kernel: bytes, n_lanes: int, args_ptr: int, global_size: tuple[int, int, int],
-                      local_size: tuple[int, int, int], program, max_steps: int, debug: bool, trace_len: int,
+                      local_size: tuple[int, int, int], max_steps: int, debug: bool, trace_len: int,
                       kernel_idx: int = 0, max_workgroups: int = 8) -> tuple[bool, str, int]:
   """Run a single kernel through both emulators. Returns (success, message, total_steps)."""
   gx, gy, gz = global_size
@@ -181,9 +186,9 @@ def run_single_kernel(kernel: bytes, n_lanes: int, args_ptr: int, global_size: t
             rust_before = rust.get_snapshot()
             python_before = python.get_snapshot()
 
-            assert python.program is not None
-            inst_info = python.program.get(python.lib_addr + python_before.pc * 4)  # Convert word offset to actual address
-            inst_hex_name = inst_info[0] if inst_info else f"unknown at PC={python_before.pc}"
+            pc_addr = python.lib_addr + python_before.pc * 4  # Convert word offset to actual address
+            python._ensure_decoded(pc_addr)
+            inst_hex_name = python.program[pc_addr][0]
             # Decode the instruction to get mnemonic for sync_after checks
             try:
               # Format is mnemonic_hexbytes, e.g. v_exp_f32_e32_014b027e -> hex is 014b027e
@@ -310,12 +315,11 @@ def compare_emulators_multi_kernel(kernels: list[KernelInfo], buf_pool: dict[int
     kernel_ranges = ranges | {(args_ptr, ctypes.sizeof(args))}
     set_valid_mem_ranges(kernel_ranges)
 
-    program = decode_program(kernel.code)
     n_lanes = kernel.local_size[0] * kernel.local_size[1] * kernel.local_size[2]
 
     ok, msg, steps = run_single_kernel(
       kernel.code, min(n_lanes, 32), args_ptr, kernel.global_size,
-      kernel.local_size, program, max_steps, debug, trace_len, ki
+      kernel.local_size, max_steps, debug, trace_len, ki
     )
     total_steps += steps
     if not ok:
@@ -341,9 +345,8 @@ def compare_emulators_with_memory(kernel: bytes, n_lanes: int, buf_sizes: list,
   ranges.add((args_ptr, ctypes.sizeof(args)))
   set_valid_mem_ranges(ranges)
 
-  program = decode_program(kernel)
   # Legacy wrapper assumes local_size = (n_lanes, 1, 1)
-  ok, msg, _ = run_single_kernel(kernel, n_lanes, args_ptr, global_size, (n_lanes, 1, 1), program, max_steps, debug, trace_len)
+  ok, msg, _ = run_single_kernel(kernel, n_lanes, args_ptr, global_size, (n_lanes, 1, 1), max_steps, debug, trace_len)
   return ok, msg
 
 def get_kernels_from_tinygrad(op_fn) -> tuple[list[KernelInfo], dict[int, int], dict[int, bytes]]:

test/amd/test_isel.py

@@ -0,0 +1,674 @@
+#!/usr/bin/env python3
+"""Tests for the pcode-based instruction selector (isel.py)."""
+import unittest
+from tinygrad.uop.ops import UOp, Ops
+from tinygrad.dtype import dtypes
+from extra.assembly.amd.isel import (rdna3_isel, make_inst, normalize, _count_nodes, _is_direct_alu,
+                                      _parse_pcode_patterns, _pattern_key, _runtime_key, uop_to_upat,
+                                      _SENTINEL, _SENTINEL_SET, _DIRECT_TABLE, _STRUCTURAL_TABLE,
+                                      _ALU_ENUM_TYPES, build_isel_patterns)
+from extra.assembly.amd.autogen.rdna3.str_pcode import PCODE
+from extra.assembly.amd.autogen.rdna3.enum import VOP2Op, VOP1Op, VOP3Op, SOP2Op, VOPCOp, VOP3SDOp
+
+# helpers
+def _var(name, dtype=dtypes.float): return UOp(Ops.DEFINE_VAR, dtype, arg=(name, 0, 100))
+def _const(val, dtype=dtypes.float): return UOp(Ops.CONST, dtype, arg=val)
+
+class TestMakeInst(unittest.TestCase):
+  def test_vop2(self):
+    inst = make_inst(VOP2Op.V_ADD_F32_E32)
+    assert inst.op == VOP2Op.V_ADD_F32_E32
+
+  def test_vop1(self):
+    inst = make_inst(VOP1Op.V_SQRT_F32_E32)
+    assert inst.op == VOP1Op.V_SQRT_F32_E32
+
+  def test_vop3(self):
+    inst = make_inst(VOP3Op.V_ADD_F64)
+    assert inst.op == VOP3Op.V_ADD_F64
+
+  def test_vop3_sdst(self):
+    # VOP3SDOp opcodes need the _SDST variant class
+    inst = make_inst(VOP3SDOp.V_ADD_CO_CI_U32)
+    assert inst.op == VOP3SDOp.V_ADD_CO_CI_U32
+
+  def test_vopc(self):
+    inst = make_inst(VOPCOp.V_CMP_LT_F32_E32)
+    assert inst.op == VOPCOp.V_CMP_LT_F32_E32
+
+  def test_sop2(self):
+    inst = make_inst(SOP2Op.S_ADD_I32)
+    assert inst.op == SOP2Op.S_ADD_I32
+
+  def test_invalid_raises(self):
+    with self.assertRaises(RuntimeError): make_inst("not_an_opcode")
+
+class TestNormalize(unittest.TestCase):
+  def test_bitcast_sentinel(self):
+    s0 = _SENTINEL['S0']
+    bc = UOp(Ops.BITCAST, dtypes.float, (s0,))
+    norm = normalize(bc)
+    assert norm.op == Ops.DEFINE_VAR
+    assert norm.dtype == dtypes.float
+
+  def test_cast_sentinel(self):
+    s0 = _SENTINEL['S0']
+    cast = UOp(Ops.CAST, dtypes.int, (s0,))
+    norm = normalize(cast)
+    assert norm.op == Ops.DEFINE_VAR
+    assert norm.dtype == dtypes.int
+
+  def test_identity_bitcast(self):
+    x = UOp(Ops.CONST, dtypes.float, arg=1.0)
+    bc = UOp(Ops.BITCAST, dtypes.float, (x,))
+    norm = normalize(bc)
+    assert norm.op == Ops.CONST
+    assert norm.arg == 1.0
+
+  def test_shift_mask_31(self):
+    s0 = _SENTINEL['S0']
+    c31 = UOp(Ops.CONST, dtypes.uint, arg=31)
+    masked = UOp(Ops.AND, dtypes.uint, (s0, c31))
+    norm = normalize(masked)
+    assert norm.op == Ops.DEFINE_VAR
+
+  def test_shift_mask_63(self):
+    s0 = _SENTINEL['S0']
+    c63 = UOp(Ops.CONST, dtypes.uint, arg=63)
+    masked = UOp(Ops.AND, dtypes.uint, (s0, c63))
+    norm = normalize(masked)
+    assert norm.op == Ops.DEFINE_VAR
+
+  def test_non_sentinel_bitcast_preserved(self):
+    x = _var('x', dtypes.uint)
+    bc = UOp(Ops.BITCAST, dtypes.float, (x,))
+    norm = normalize(bc)
+    assert norm.op == Ops.BITCAST  # not a sentinel, so preserved
+
+  def test_recursive(self):
+    s0 = _SENTINEL['S0']
+    s1 = _SENTINEL['S1']
+    bc0 = UOp(Ops.BITCAST, dtypes.float, (s0,))
+    bc1 = UOp(Ops.BITCAST, dtypes.float, (s1,))
+    add = UOp(Ops.ADD, dtypes.float, (bc0, bc1))
+    norm = normalize(add)
+    assert norm.op == Ops.ADD
+    assert all(s.dtype == dtypes.float for s in norm.src)
+    assert all(s.op == Ops.DEFINE_VAR for s in norm.src)
+
+class TestCountNodes(unittest.TestCase):
+  def test_leaf(self):
+    assert _count_nodes(_var('x')) == 1
+
+  def test_binary(self):
+    x, y = _var('x'), _var('y')
+    add = UOp(Ops.ADD, dtypes.float, (x, y))
+    assert _count_nodes(add) == 3
+
+  def test_dag_sharing(self):
+    x = _var('x')
+    add = UOp(Ops.ADD, dtypes.float, (x, x))
+    assert _count_nodes(add) == 2  # x counted once
+
+class TestIsDirectAlu(unittest.TestCase):
+  def test_add_sentinels(self):
+    s0 = _SENTINEL['S0'].replace(dtype=dtypes.float)
+    s1 = _SENTINEL['S1'].replace(dtype=dtypes.float)
+    add = UOp(Ops.ADD, dtypes.float, (s0, s1))
+    assert _is_direct_alu(add)
+
+  def test_cast_sentinel(self):
+    s0 = _SENTINEL['S0'].replace(dtype=dtypes.int)
+    cast = UOp(Ops.CAST, dtypes.float, (s0,))
+    assert _is_direct_alu(cast)
+
+  def test_nested_not_direct(self):
+    s0 = _SENTINEL['S0'].replace(dtype=dtypes.uint)
+    c = _const(0xFFFFFFFF, dtypes.uint)
+    xor = UOp(Ops.XOR, dtypes.uint, (s0, c))
+    assert not _is_direct_alu(xor)  # const child is not DEFINE_VAR
+
+class TestPatternKey(unittest.TestCase):
+  def test_sentinel_var(self):
+    s0 = _SENTINEL['S0'].replace(dtype=dtypes.float)
+    key = _pattern_key(s0)
+    assert key == 'var(S0,dtypes.float)'
+
+  def test_const(self):
+    c = _const(42, dtypes.uint)
+    key = _pattern_key(c)
+    assert key == 'const(42,dtypes.uint)'
+
+  def test_binary_op(self):
+    s0 = _SENTINEL['S0'].replace(dtype=dtypes.float)
+    s1 = _SENTINEL['S1'].replace(dtype=dtypes.float)
+    add = UOp(Ops.ADD, dtypes.float, (s0, s1))
+    key = _pattern_key(add)
+    assert key == 'Ops.ADD(dtypes.float,var(S0,dtypes.float),var(S1,dtypes.float))'
+
+class TestRuntimeKey(unittest.TestCase):
+  def test_matches_pattern_key(self):
+    # runtime key on a matched UOp should equal pattern key on the pcode template
+    x = _var('x', dtypes.uint)
+    c = _const(0xFFFFFFFF, dtypes.uint)
+    xor = UOp(Ops.XOR, dtypes.uint, (x, c))
+    rkey = _runtime_key(xor)
+
+    s0 = _SENTINEL['S0'].replace(dtype=dtypes.uint)
+    xor_template = UOp(Ops.XOR, dtypes.uint, (s0, c))
+    pkey = _pattern_key(xor_template)
+    assert rkey == pkey
+
+class TestUopToUpat(unittest.TestCase):
+  def test_sentinel_becomes_var(self):
+    s0 = _SENTINEL['S0'].replace(dtype=dtypes.float)
+    pat = uop_to_upat(s0)
+    assert pat.name == 'S0'
+    assert pat.dtype == (dtypes.float,)
+
+  def test_const_preserved(self):
+    c = _const(42, dtypes.uint)
+    pat = uop_to_upat(c)
+    assert pat.op == (Ops.CONST,)
+    assert pat.arg == 42
+
+class TestBuildPerformance(unittest.TestCase):
+  def test_builds_under_2_seconds(self):
+    import time
+    t0 = time.time()
+    build_isel_patterns(PCODE)
+    elapsed = time.time() - t0
+    assert elapsed < 2.0, f"build took {elapsed:.2f}s, expected <2s"
+
+  def test_alu_filter(self):
+    # verify only ALU enum types are parsed
+    for opcode in PCODE:
+      if type(opcode).__name__ not in _ALU_ENUM_TYPES: continue
+      # these should parse without hanging
+
+class TestDirectPatterns(unittest.TestCase):
+  @classmethod
+  def setUpClass(cls):
+    cls.pm = rdna3_isel()
+
+  def _check(self, uop, expected_name_substr):
+    result = self.pm.rewrite(uop)
+    self.assertIsNotNone(result, f"no match for {uop.op} {uop.dtype}")
+    self.assertEqual(result.op, Ops.INS)
+    self.assertIn(expected_name_substr, result.arg.op.name, f"expected {expected_name_substr} in {result.arg.op.name}")
+    return result
+
+  # arithmetic
+  def test_add_f32(self): self._check(UOp(Ops.ADD, dtypes.float, (_var('a'), _var('b'))), 'V_ADD_F32')
+  def test_add_f64(self): self._check(UOp(Ops.ADD, dtypes.double, (_var('a', dtypes.double), _var('b', dtypes.double))), 'V_ADD_F64')
+  def test_add_i32(self): self._check(UOp(Ops.ADD, dtypes.int, (_var('a', dtypes.int), _var('b', dtypes.int))), 'ADD_NC_I32')
+  def test_add_u32(self): self._check(UOp(Ops.ADD, dtypes.uint, (_var('a', dtypes.uint), _var('b', dtypes.uint))), 'ADD_NC_U32')
+  def test_mul_f32(self): self._check(UOp(Ops.MUL, dtypes.float, (_var('a'), _var('b'))), 'V_MUL_F32')
+  def test_mul_f64(self): self._check(UOp(Ops.MUL, dtypes.double, (_var('a', dtypes.double), _var('b', dtypes.double))), 'V_MUL_F64')
+  def test_mul_i32(self): self._check(UOp(Ops.MUL, dtypes.int, (_var('a', dtypes.int), _var('b', dtypes.int))), 'MUL_I32')
+  def test_mul_u32(self): self._check(UOp(Ops.MUL, dtypes.uint, (_var('a', dtypes.uint), _var('b', dtypes.uint))), 'MUL_U32')
+
+  # bitwise
+  def test_and_u32(self): self._check(UOp(Ops.AND, dtypes.uint, (_var('a', dtypes.uint), _var('b', dtypes.uint))), 'AND_B32')
+  def test_or_u32(self): self._check(UOp(Ops.OR, dtypes.uint, (_var('a', dtypes.uint), _var('b', dtypes.uint))), 'OR_B32')
+  def test_xor_u32(self): self._check(UOp(Ops.XOR, dtypes.uint, (_var('a', dtypes.uint), _var('b', dtypes.uint))), 'XOR_B32')
+  # u64 bitwise ops are SOP-only, skipped in vgpr_only mode
+  def test_and_u64_skipped(self):
+    result = self.pm.rewrite(UOp(Ops.AND, dtypes.ulong, (_var('a', dtypes.ulong), _var('b', dtypes.ulong))))
+    self.assertIsNone(result)
+  def test_or_u64_skipped(self):
+    result = self.pm.rewrite(UOp(Ops.OR, dtypes.ulong, (_var('a', dtypes.ulong), _var('b', dtypes.ulong))))
+    self.assertIsNone(result)
+  def test_xor_u64_skipped(self):
+    result = self.pm.rewrite(UOp(Ops.XOR, dtypes.ulong, (_var('a', dtypes.ulong), _var('b', dtypes.ulong))))
+    self.assertIsNone(result)
+
+  # shifts
+  def test_shl_u32(self): self._check(UOp(Ops.SHL, dtypes.uint, (_var('a', dtypes.uint), _var('b', dtypes.uint))), 'LSH')
+  def test_shr_u32(self): self._check(UOp(Ops.SHR, dtypes.uint, (_var('a', dtypes.uint), _var('b', dtypes.uint))), 'LSH')
+
+  # unary float
+  def test_sqrt_f32(self): self._check(UOp(Ops.SQRT, dtypes.float, (_var('a'),)), 'SQRT_F32')
+  def test_sqrt_f64(self): self._check(UOp(Ops.SQRT, dtypes.double, (_var('a', dtypes.double),)), 'SQRT_F64')
+  def test_trunc_f32(self): self._check(UOp(Ops.TRUNC, dtypes.float, (_var('a'),)), 'TRUNC_F32')
+  def test_trunc_f64(self): self._check(UOp(Ops.TRUNC, dtypes.double, (_var('a', dtypes.double),)), 'TRUNC_F64')
+  def test_log2_f32(self): self._check(UOp(Ops.LOG2, dtypes.float, (_var('a'),)), 'LOG')
+  def test_exp2_f32(self): self._check(UOp(Ops.EXP2, dtypes.float, (_var('a'),)), 'EXP')
+
+  # conversions
+  def test_cast_i32_to_f32(self): self._check(UOp(Ops.CAST, dtypes.float, (_var('a', dtypes.int),)), 'CVT_F32_I32')
+  def test_cast_f32_to_f64(self): self._check(UOp(Ops.CAST, dtypes.double, (_var('a'),)), 'CVT_F64_F32')
+  def test_cast_f64_to_f32(self): self._check(UOp(Ops.CAST, dtypes.float, (_var('a', dtypes.double),)), 'CVT_F32_F64')
+  def test_cast_i32_to_f64(self): self._check(UOp(Ops.CAST, dtypes.double, (_var('a', dtypes.int),)), 'CVT_F64_I32')
+  def test_cast_f32_to_f16(self): self._check(UOp(Ops.CAST, dtypes.half, (_var('a'),)), 'CVT_F16_F32')
+
+  # compares are skipped by ISel (VOPC writes VCC, not VGPRs; LLVM handles natively)
+  def test_cmplt_skipped(self):
+    result = self.pm.rewrite(UOp(Ops.CMPLT, dtypes.bool, (_var('a', dtypes.int), _var('b', dtypes.int))))
+    self.assertIsNone(result)
+  def test_cmpne_skipped(self):
+    result = self.pm.rewrite(UOp(Ops.CMPNE, dtypes.bool, (_var('a'), _var('b'))))
+    self.assertIsNone(result)
+
+  # check that unmatched types return None
+  def test_no_match(self):
+    # there's no direct ADD for bools
+    result = self.pm.rewrite(UOp(Ops.ADD, dtypes.bool, (_var('a', dtypes.bool), _var('b', dtypes.bool))))
+    self.assertIsNone(result)
+
+class TestStructuralPatterns(unittest.TestCase):
+  @classmethod
+  def setUpClass(cls):
+    cls.pm = rdna3_isel()
+
+  def _check(self, uop, expected_name_substr, expected_src_count=None):
+    result = self.pm.rewrite(uop)
+    self.assertIsNotNone(result, f"no match for structural pattern")
+    self.assertEqual(result.op, Ops.INS)
+    self.assertIn(expected_name_substr, result.arg.op.name, f"expected {expected_name_substr} in {result.arg.op.name}")
+    if expected_src_count is not None:
+      self.assertEqual(len(result.src), expected_src_count, f"expected {expected_src_count} srcs, got {len(result.src)}")
+    return result
+
+  def test_not_u32(self):
+    x = _var('x', dtypes.uint)
+    xor = UOp(Ops.XOR, dtypes.uint, (x, _const(0xFFFFFFFF, dtypes.uint)))
+    self._check(xor, 'NOT_B32', 1)
+
+  # u64 NOT is SOP-only (S_NOT_B64), skipped in vgpr_only mode
+  def test_not_u64_skipped(self):
+    x = _var('x', dtypes.ulong)
+    xor = UOp(Ops.XOR, dtypes.ulong, (x, _const(0xFFFFFFFFFFFFFFFF, dtypes.ulong)))
+    result = self.pm.rewrite(xor)
+    self.assertIsNone(result)
+
+  def test_sub_u32(self):
+    x, y = _var('x', dtypes.uint), _var('y', dtypes.uint)
+    neg = UOp(Ops.MUL, dtypes.uint, (y, _const(-1, dtypes.uint)))
+    sub = UOp(Ops.ADD, dtypes.uint, (x, neg))
+    result = self._check(sub, 'SUB_NC_U32', 2)
+    # verify source order: x is first, y is second
+    self.assertEqual(result.src[0].arg, ('x', 0, 100))
+    self.assertEqual(result.src[1].arg, ('y', 0, 100))
+
+  def test_rcp_f32(self):
+    a = _var('a')
+    rcp = UOp(Ops.RECIPROCAL, dtypes.float, (a,))
+    mul_rcp = UOp(Ops.MUL, dtypes.float, (_const(1.0), rcp))
+    result = self._check(mul_rcp, 'RCP_F32', 1)
+    self.assertEqual(result.src[0].arg, ('a', 0, 100))
+
+  def test_rcp_f64(self):
+    a = _var('a', dtypes.double)
+    rcp = UOp(Ops.RECIPROCAL, dtypes.double, (a,))
+    mul_rcp = UOp(Ops.MUL, dtypes.double, (_const(1.0, dtypes.double), rcp))
+    self._check(mul_rcp, 'RCP_F64', 1)
+
+  def test_cvt_i32_f32(self):
+    # CAST(i32, TRUNC(f32, x)) -> V_CVT_I32_F32
+    a = _var('a')
+    trunc = UOp(Ops.TRUNC, dtypes.float, (a,))
+    cast = UOp(Ops.CAST, dtypes.int, (trunc,))
+    self._check(cast, 'CVT_I32_F32', 1)
+
+  def test_mad_u32(self):
+    x, y, z = _var('x', dtypes.uint), _var('y', dtypes.uint), _var('z', dtypes.uint)
+    mul = UOp(Ops.MUL, dtypes.uint, (x, y))
+    mad = UOp(Ops.ADD, dtypes.uint, (mul, z))
+    result = self._check(mad, 'MAD_U32_U24', 3)
+    self.assertEqual(result.src[0].arg, ('x', 0, 100))
+    self.assertEqual(result.src[1].arg, ('y', 0, 100))
+    self.assertEqual(result.src[2].arg, ('z', 0, 100))
+
+  def test_add3_u32(self):
+    x, y, z = _var('x', dtypes.uint), _var('y', dtypes.uint), _var('z', dtypes.uint)
+    add1 = UOp(Ops.ADD, dtypes.uint, (x, y))
+    add3 = UOp(Ops.ADD, dtypes.uint, (add1, z))
+    result = self._check(add3, 'ADD3_U32', 3)
+
+  def test_xor3_b32(self):
+    x, y, z = _var('x', dtypes.uint), _var('y', dtypes.uint), _var('z', dtypes.uint)
+    xor1 = UOp(Ops.XOR, dtypes.uint, (x, y))
+    xor3 = UOp(Ops.XOR, dtypes.uint, (xor1, z))
+    self._check(xor3, 'XOR3_B32', 3)
+
+  def test_and_or_b32(self):
+    x, y, z = _var('x', dtypes.uint), _var('y', dtypes.uint), _var('z', dtypes.uint)
+    and_op = UOp(Ops.AND, dtypes.uint, (x, y))
+    or_op = UOp(Ops.OR, dtypes.uint, (and_op, z))
+    self._check(or_op, 'AND_OR_B32', 3)
+
+  def test_or3_b32(self):
+    x, y, z = _var('x', dtypes.uint), _var('y', dtypes.uint), _var('z', dtypes.uint)
+    or1 = UOp(Ops.OR, dtypes.uint, (x, y))
+    or3 = UOp(Ops.OR, dtypes.uint, (or1, z))
+    self._check(or3, 'OR3_B32', 3)
+
+  # NAND/NOR were SOP-only, in vgpr_only mode they decompose to V_XOR_B32(AND/OR, mask)
+  def test_nand_decomposes(self):
+    x, y = _var('x', dtypes.uint), _var('y', dtypes.uint)
+    and_op = UOp(Ops.AND, dtypes.uint, (x, y))
+    nand = UOp(Ops.XOR, dtypes.uint, (and_op, _const(0xFFFFFFFF, dtypes.uint)))
+    self._check(nand, 'XOR_B32')
+
+  def test_nor_decomposes(self):
+    x, y = _var('x', dtypes.uint), _var('y', dtypes.uint)
+    or_op = UOp(Ops.OR, dtypes.uint, (x, y))
+    nor = UOp(Ops.XOR, dtypes.uint, (or_op, _const(0xFFFFFFFF, dtypes.uint)))
+    self._check(nor, 'XOR_B32')
+
+  def test_xnor_b32(self):
+    x, y = _var('x', dtypes.uint), _var('y', dtypes.uint)
+    xor_op = UOp(Ops.XOR, dtypes.uint, (x, y))
+    xnor = UOp(Ops.XOR, dtypes.uint, (xor_op, _const(0xFFFFFFFF, dtypes.uint)))
+    self._check(xnor, 'XNOR_B32', 2)
+
+  def test_min_u32(self):
+    x, y = _var('x', dtypes.uint), _var('y', dtypes.uint)
+    cmp = UOp(Ops.CMPLT, dtypes.bool, (x, y))
+    where = UOp(Ops.WHERE, dtypes.uint, (cmp, x, y))
+    self._check(where, 'MIN_U32', 2)
+
+class TestInstProperties(unittest.TestCase):
+  """Verify that Inst objects produced by isel have correct properties."""
+  @classmethod
+  def setUpClass(cls):
+    cls.pm = rdna3_isel()
+
+  def test_ins_has_dtype(self):
+    result = self.pm.rewrite(UOp(Ops.ADD, dtypes.float, (_var('a'), _var('b'))))
+    self.assertEqual(result.dtype, dtypes.float)
+
+  def test_ins_preserves_sources(self):
+    a, b = _var('a'), _var('b')
+    result = self.pm.rewrite(UOp(Ops.ADD, dtypes.float, (a, b)))
+    self.assertEqual(result.src, (a, b))
+
+  def test_ins_tag_default_none(self):
+    result = self.pm.rewrite(UOp(Ops.ADD, dtypes.float, (_var('a'), _var('b'))))
+    # tag should not be set (defaults to None or empty)
+    self.assertIsNone(result.tag)
+
+class TestTableCoverage(unittest.TestCase):
+  """Verify that the tables have expected coverage."""
+  @classmethod
+  def setUpClass(cls):
+    rdna3_isel()  # populate tables
+
+  def test_direct_table_has_add(self):
+    found = any(op == Ops.ADD for (op, _, _) in _DIRECT_TABLE)
+    self.assertTrue(found)
+
+  def test_direct_table_has_cast(self):
+    found = any(op == Ops.CAST for (op, _, _) in _DIRECT_TABLE)
+    self.assertTrue(found)
+
+  def test_direct_table_skips_cmplt(self):
+    # compares are in the table but skipped at runtime (bool output)
+    found = any(op == Ops.CMPLT for (op, _, _) in _DIRECT_TABLE)
+    self.assertTrue(found)  # entries exist but callbacks skip them
+
+  def test_structural_table_has_not(self):
+    found = any('NOT' in inst.op.name for inst in _STRUCTURAL_TABLE.values())
+    self.assertTrue(found)
+
+  def test_structural_table_has_rcp(self):
+    found = any('RCP' in inst.op.name for inst in _STRUCTURAL_TABLE.values())
+    self.assertTrue(found)
+
+  def test_structural_table_has_sub(self):
+    found = any('SUB' in inst.op.name for inst in _STRUCTURAL_TABLE.values())
+    self.assertTrue(found)
+
+  def test_direct_count(self):
+    self.assertGreaterEqual(len(_DIRECT_TABLE), 25, "expected at least 25 direct patterns")
+
+  def test_structural_count(self):
+    self.assertGreaterEqual(len(_STRUCTURAL_TABLE), 15, "expected at least 15 structural patterns")
+
+class TestEmulatorValidation(unittest.TestCase):
+  """Validate isel-produced Inst objects execute correctly in the emulator.
+
+  For each pattern, we:
+  1. Run the UOp through isel to get Ops.INS with arg=Inst
+  2. Copy the Inst and assign concrete registers
+  3. Set up operand values via MOV instructions
+  4. Execute through the emulator
+  5. Verify the output matches expected computation
+  """
+  @classmethod
+  def setUpClass(cls):
+    cls.pm = rdna3_isel()
+
+  def _run(self, instructions, n_lanes=1):
+    from extra.assembly.amd.test.hw.helpers import run_program_emu
+    return run_program_emu(instructions, n_lanes)
+
+  def _get_inst(self, uop):
+    """Get isel result, return (Inst, src_count)."""
+    result = self.pm.rewrite(uop)
+    assert result is not None and result.op == Ops.INS, f"isel failed for {uop.op} {uop.dtype}"
+    return result.arg, len(result.src)
+
+  def _copy_inst(self, inst):
+    import copy
+    return copy.copy(inst)
+
+  # ── direct ALU: float arithmetic ──
+
+  def test_emu_add_f32(self):
+    from extra.assembly.amd.test.hw.helpers import i2f, f2i
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    inst, _ = self._get_inst(UOp(Ops.ADD, dtypes.float, (_var('a'), _var('b'))))
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.vsrc1 = v[1]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 1.5), v_mov_b32_e32(v[1], 2.25), ci])
+    self.assertAlmostEqual(i2f(st.vgpr[0][2]), 3.75, places=5)
+
+  def test_emu_mul_f32(self):
+    from extra.assembly.amd.test.hw.helpers import i2f
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    inst, _ = self._get_inst(UOp(Ops.MUL, dtypes.float, (_var('a'), _var('b'))))
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.vsrc1 = v[1]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 3.0), v_mov_b32_e32(v[1], 4.0), ci])
+    self.assertAlmostEqual(i2f(st.vgpr[0][2]), 12.0, places=5)
+
+  # ── direct ALU: integer arithmetic ──
+
+  def test_emu_add_u32(self):
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    inst, _ = self._get_inst(UOp(Ops.ADD, dtypes.uint, (_var('a', dtypes.uint), _var('b', dtypes.uint))))
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.vsrc1 = v[1]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 10), v_mov_b32_e32(v[1], 20), ci])
+    self.assertEqual(st.vgpr[0][2], 30)
+
+  # ── direct ALU: bitwise ──
+
+  def test_emu_and_u32(self):
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    inst, _ = self._get_inst(UOp(Ops.AND, dtypes.uint, (_var('a', dtypes.uint), _var('b', dtypes.uint))))
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.vsrc1 = v[1]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 0xFF00), v_mov_b32_e32(v[1], 0x0FF0), ci])
+    self.assertEqual(st.vgpr[0][2], 0x0F00)
+
+  def test_emu_or_u32(self):
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    inst, _ = self._get_inst(UOp(Ops.OR, dtypes.uint, (_var('a', dtypes.uint), _var('b', dtypes.uint))))
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.vsrc1 = v[1]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 0xFF00), v_mov_b32_e32(v[1], 0x0FF0), ci])
+    self.assertEqual(st.vgpr[0][2], 0xFFF0)
+
+  def test_emu_xor_u32(self):
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    inst, _ = self._get_inst(UOp(Ops.XOR, dtypes.uint, (_var('a', dtypes.uint), _var('b', dtypes.uint))))
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.vsrc1 = v[1]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 0xFF00), v_mov_b32_e32(v[1], 0x0FF0), ci])
+    self.assertEqual(st.vgpr[0][2], 0xF0F0)
+
+  # ── direct ALU: shifts ──
+
+  def test_emu_shl_u32(self):
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    inst, _ = self._get_inst(UOp(Ops.SHL, dtypes.uint, (_var('a', dtypes.uint), _var('b', dtypes.uint))))
+    ci = self._copy_inst(inst)
+    # LSHLREV: vdst = vsrc1 << src0 (reversed operands!)
+    ci.src0 = v[1]; ci.vsrc1 = v[0]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 1), v_mov_b32_e32(v[1], 4), ci])
+    self.assertEqual(st.vgpr[0][2], 16)  # 1 << 4 = 16
+
+  def test_emu_shr_u32(self):
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    inst, _ = self._get_inst(UOp(Ops.SHR, dtypes.uint, (_var('a', dtypes.uint), _var('b', dtypes.uint))))
+    ci = self._copy_inst(inst)
+    # LSHRREV: vdst = vsrc1 >> src0 (reversed operands!)
+    ci.src0 = v[1]; ci.vsrc1 = v[0]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 16), v_mov_b32_e32(v[1], 4), ci])
+    self.assertEqual(st.vgpr[0][2], 1)  # 16 >> 4 = 1
+
+  # ── direct ALU: unary float ──
+
+  def test_emu_sqrt_f32(self):
+    from extra.assembly.amd.test.hw.helpers import i2f
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    inst, _ = self._get_inst(UOp(Ops.SQRT, dtypes.float, (_var('a'),)))
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 4.0), ci])
+    self.assertAlmostEqual(i2f(st.vgpr[0][2]), 2.0, places=4)
+
+  def test_emu_trunc_f32(self):
+    from extra.assembly.amd.test.hw.helpers import i2f
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    inst, _ = self._get_inst(UOp(Ops.TRUNC, dtypes.float, (_var('a'),)))
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 3.7), ci])
+    self.assertAlmostEqual(i2f(st.vgpr[0][2]), 3.0, places=5)
+
+  def test_emu_exp2_f32(self):
+    from extra.assembly.amd.test.hw.helpers import i2f
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    inst, _ = self._get_inst(UOp(Ops.EXP2, dtypes.float, (_var('a'),)))
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 3.0), ci])
+    self.assertAlmostEqual(i2f(st.vgpr[0][2]), 8.0, delta=0.01)
+
+  def test_emu_log2_f32(self):
+    from extra.assembly.amd.test.hw.helpers import i2f
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    inst, _ = self._get_inst(UOp(Ops.LOG2, dtypes.float, (_var('a'),)))
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 8.0), ci])
+    self.assertAlmostEqual(i2f(st.vgpr[0][2]), 3.0, delta=0.01)
+
+  # ── direct ALU: conversions ──
+
+  def test_emu_cast_i32_to_f32(self):
+    from extra.assembly.amd.test.hw.helpers import i2f
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    inst, _ = self._get_inst(UOp(Ops.CAST, dtypes.float, (_var('a', dtypes.int),)))
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 42), ci])
+    self.assertAlmostEqual(i2f(st.vgpr[0][2]), 42.0, places=5)
+
+  def test_emu_cast_f32_to_f16(self):
+    from extra.assembly.amd.test.hw.helpers import i2f, f16
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    inst, _ = self._get_inst(UOp(Ops.CAST, dtypes.half, (_var('a'),)))
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 1.5), ci])
+    # f16 result is in lower 16 bits of v[2]
+    self.assertAlmostEqual(f16(st.vgpr[0][2]), 1.5, places=2)
+
+  # compares are skipped by ISel (VOPC writes VCC; LLVM handles natively)
+
+  # ── structural: NOT ──
+
+  def test_emu_not_u32(self):
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    x = _var('x', dtypes.uint)
+    xor = UOp(Ops.XOR, dtypes.uint, (x, _const(0xFFFFFFFF, dtypes.uint)))
+    inst, _ = self._get_inst(xor)
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 0x0000FF00), ci])
+    self.assertEqual(st.vgpr[0][2], 0xFFFF00FF)
+
+  # ── structural: SUB ──
+
+  def test_emu_sub_u32(self):
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    x, y = _var('x', dtypes.uint), _var('y', dtypes.uint)
+    neg = UOp(Ops.MUL, dtypes.uint, (y, _const(-1, dtypes.uint)))
+    sub = UOp(Ops.ADD, dtypes.uint, (x, neg))
+    inst, nsrc = self._get_inst(sub)
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.vsrc1 = v[1]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 30), v_mov_b32_e32(v[1], 12), ci])
+    self.assertEqual(st.vgpr[0][2], 18)
+
+  # ── structural: RCP ──
+
+  def test_emu_rcp_f32(self):
+    from extra.assembly.amd.test.hw.helpers import i2f
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    a = _var('a')
+    rcp = UOp(Ops.RECIPROCAL, dtypes.float, (a,))
+    mul_rcp = UOp(Ops.MUL, dtypes.float, (_const(1.0), rcp))
+    inst, _ = self._get_inst(mul_rcp)
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.vdst = v[2]
+    st = self._run([v_mov_b32_e32(v[0], 4.0), ci])
+    self.assertAlmostEqual(i2f(st.vgpr[0][2]), 0.25, places=4)
+
+  # ── structural: MAD ──
+
+  def test_emu_mad_u32(self):
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    x, y, z = _var('x', dtypes.uint), _var('y', dtypes.uint), _var('z', dtypes.uint)
+    mul = UOp(Ops.MUL, dtypes.uint, (x, y))
+    mad = UOp(Ops.ADD, dtypes.uint, (mul, z))
+    inst, _ = self._get_inst(mad)
+    ci = self._copy_inst(inst)
+    # VOP3 format: src0, src1, src2, vdst
+    ci.src0 = v[0]; ci.src1 = v[1]; ci.src2 = v[2]; ci.vdst = v[3]
+    st = self._run([v_mov_b32_e32(v[0], 3), v_mov_b32_e32(v[1], 4), v_mov_b32_e32(v[2], 5), ci])
+    self.assertEqual(st.vgpr[0][3], 17)  # 3*4 + 5 = 17
+
+  # ── structural: ADD3 ──
+
+  def test_emu_add3_u32(self):
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    x, y, z = _var('x', dtypes.uint), _var('y', dtypes.uint), _var('z', dtypes.uint)
+    add1 = UOp(Ops.ADD, dtypes.uint, (x, y))
+    add3 = UOp(Ops.ADD, dtypes.uint, (add1, z))
+    inst, _ = self._get_inst(add3)
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.src1 = v[1]; ci.src2 = v[2]; ci.vdst = v[3]
+    st = self._run([v_mov_b32_e32(v[0], 10), v_mov_b32_e32(v[1], 20), v_mov_b32_e32(v[2], 30), ci])
+    self.assertEqual(st.vgpr[0][3], 60)  # 10+20+30
+
+  # ── structural: XOR3 ──
+
+  def test_emu_xor3_b32(self):
+    from extra.assembly.amd.autogen.rdna3.ins import v, v_mov_b32_e32
+    x, y, z = _var('x', dtypes.uint), _var('y', dtypes.uint), _var('z', dtypes.uint)
+    xor1 = UOp(Ops.XOR, dtypes.uint, (x, y))
+    xor3 = UOp(Ops.XOR, dtypes.uint, (xor1, z))
+    inst, _ = self._get_inst(xor3)
+    ci = self._copy_inst(inst)
+    ci.src0 = v[0]; ci.src1 = v[1]; ci.src2 = v[2]; ci.vdst = v[3]
+    st = self._run([v_mov_b32_e32(v[0], 0xFF), v_mov_b32_e32(v[1], 0x0F), v_mov_b32_e32(v[2], 0x33), ci])
+    self.assertEqual(st.vgpr[0][3], 0xFF ^ 0x0F ^ 0x33)
+
+if __name__ == '__main__':
+  unittest.main()

test/amd/test_rdna4_emu.py

@@ -1,96 +0,0 @@
-import unittest, ctypes
-from tinygrad.runtime.autogen.amd.rdna4 import ins as ir4
-from tinygrad.renderer.amd.dsl import v, s
-from tinygrad.renderer.amd.emu import WaveState, decode_program
-from tinygrad.device import Buffer, BufferSpec
-from tinygrad.dtype import dtypes
-
-class TestRDNA4Emu(unittest.TestCase):
-  def _run(self, insts: list, sgprs: dict[int, int] | None = None, vgprs: dict[tuple[int, int], int] | None = None) -> WaveState:
-    """Run instructions and return final WaveState."""
-    # Add S_ENDPGM if not present
-    if not any(isinstance(i, ir4.SOPP) and i.op == ir4.SOPPOp.S_ENDPGM for i in insts):
-      insts = list(insts) + [ir4.SOPP(ir4.SOPPOp.S_ENDPGM, simm=0)]
-
-    # Assemble and decode
-    code = b''.join(i.to_bytes() for i in insts)
-    code_buf = (ctypes.c_uint8 * len(code)).from_buffer_copy(code)
-    code_addr = ctypes.addressof(code_buf)
-    program_raw = decode_program(code, "rdna4")
-    program = {code_addr + offset: val for offset, val in program_raw.items()}
-
-    # Setup wave state
-    st = WaveState(n_lanes=1)
-    st.pc = code_addr
-    for idx, val in (sgprs or {}).items(): st._write_sgpr(idx, val)
-    for (reg, lane), val in (vgprs or {}).items(): st._write_vgpr(reg, lane, val)
-
-    # Setup vmem buffer with external_ptr=0 (maps to address 0, allows any pointer access)
-    vmem_buf = Buffer('CPU', 1 << 40, dtypes.uint32, options=BufferSpec(external_ptr=0)).ensure_allocated()
-
-    # Execute
-    c_bufs = [ctypes.c_uint64(st.sgpr_buf._buf.va_addr), ctypes.c_uint64(st.vgpr_buf._buf.va_addr),
-              ctypes.c_uint64(vmem_buf._buf.va_addr), ctypes.c_uint64(0), ctypes.c_uint64(0)]
-    for _ in range(100):
-      if (pc := st.pc) == 0xFFFFFFFFFFFFFFFF or pc not in program: break
-      _, fxn, globals_list, _ = program[pc]
-      fxn(*[c_bufs[g] for g in globals_list])
-    return st
-
-  def test_vopd_dual_mov(self):
-    """Test VOPD with two V_DUAL_MOV_B32 operations: v[1]=s[1], v[2]=s[2]."""
-    insts = [ir4.VOPD(ir4.VOPDOp.V_DUAL_MOV_B32, ir4.VOPDOp.V_DUAL_MOV_B32,
-                      vdstx=v[1], vdsty=v[2], srcx0=s[1], srcy0=s[2], vsrcx1=v[0], vsrcy1=v[0])]
-    st = self._run(insts, sgprs={1: 0x40e00000, 2: 0x41100000})  # 7.0f, 9.0f
-    self.assertEqual(st._read_vgpr(1, 0), 0x40e00000)  # v[1] = 7.0
-    self.assertEqual(st._read_vgpr(2, 0), 0x41100000)  # v[2] = 9.0
-
-  def test_vopd_dual_mov_after_other_vopd(self):
-    """Test VOPD reuse: first VOPD(v[3]=0, v[0]=?), then VOPD(v[1]=s[1], v[2]=s[2])."""
-    # This matches the BEAM kernel sequence that fails
-    insts = [
-      ir4.VOPD(ir4.VOPDOp.V_DUAL_MOV_B32, ir4.VOPDOp.V_DUAL_MOV_B32,
-               vdstx=v[3], vdsty=v[0], srcx0=0, srcy0=s[0], vsrcx1=v[0], vsrcy1=v[0]),  # v[3]=0, v[0]=s[0]
-      ir4.VOPD(ir4.VOPDOp.V_DUAL_MOV_B32, ir4.VOPDOp.V_DUAL_MOV_B32,
-               vdstx=v[1], vdsty=v[2], srcx0=s[1], srcy0=s[2], vsrcx1=v[0], vsrcy1=v[0]),  # v[1]=s[1], v[2]=s[2]
-    ]
-    st = self._run(insts, sgprs={0: 0x40a00000, 1: 0x40e00000, 2: 0x41100000})  # 5.0f, 7.0f, 9.0f
-    self.assertEqual(st._read_vgpr(1, 0), 0x40e00000)  # v[1] = 7.0
-    self.assertEqual(st._read_vgpr(2, 0), 0x41100000)  # v[2] = 9.0
-
-  def test_vopd_with_s_add_f32_sequence(self):
-    """Test full BEAM kernel sequence: s_add_f32 then VOPD."""
-    # This is the exact sequence from the failing BEAM kernel
-    insts = [
-      ir4.SOP2(ir4.SOP2Op.S_ADD_F32, sdst=s[0], ssrc0=s[0], ssrc1=s[8]),   # s[0] = s[0] + s[8]
-      ir4.SOP2(ir4.SOP2Op.S_ADD_F32, sdst=s[1], ssrc0=s[1], ssrc1=s[9]),   # s[1] = s[1] + s[9]
-      ir4.SOP2(ir4.SOP2Op.S_ADD_F32, sdst=s[2], ssrc0=s[2], ssrc1=s[10]),  # s[2] = s[2] + s[10]
-      ir4.VOPD(ir4.VOPDOp.V_DUAL_MOV_B32, ir4.VOPDOp.V_DUAL_MOV_B32,
-               vdstx=v[3], vdsty=v[0], srcx0=0, srcy0=s[0], vsrcx1=v[0], vsrcy1=v[0]),
-      ir4.VOPD(ir4.VOPDOp.V_DUAL_MOV_B32, ir4.VOPDOp.V_DUAL_MOV_B32,
-               vdstx=v[1], vdsty=v[2], srcx0=s[1], srcy0=s[2], vsrcx1=v[0], vsrcy1=v[0]),
-    ]
-    # Input: s[0:2] = [1,2,3], s[8:10] = [4,5,6]
-    # After s_add_f32: s[0:2] = [5,7,9]
-    st = self._run(insts, sgprs={0: 0x3f800000, 1: 0x40000000, 2: 0x40400000,  # 1.0, 2.0, 3.0
-                                  8: 0x40800000, 9: 0x40a00000, 10: 0x40c00000})  # 4.0, 5.0, 6.0
-    self.assertEqual(st._read_vgpr(1, 0), 0x40e00000)  # v[1] = 7.0
-    self.assertEqual(st._read_vgpr(2, 0), 0x41100000)  # v[2] = 9.0
-
-  def test_s_mov_b32_then_vopd(self):
-    """Test s_mov_b32 followed by VOPD - simulates BEAM kernel sequence."""
-    # Use s_mov_b32 with SGPR source (copy from pre-initialized SGPRs)
-    # s[10:12] will have values set by test harness, copy to s[0:2], then VOPD to VGPRs
-    insts = [
-      ir4.SOP1(ir4.SOP1Op.S_MOV_B32, sdst=s[0], ssrc0=s[10]),  # s[0] = s[10]
-      ir4.SOP1(ir4.SOP1Op.S_MOV_B32, sdst=s[1], ssrc0=s[11]),  # s[1] = s[11]
-      ir4.SOP1(ir4.SOP1Op.S_MOV_B32, sdst=s[2], ssrc0=s[12]),  # s[2] = s[12]
-      ir4.VOPD(ir4.VOPDOp.V_DUAL_MOV_B32, ir4.VOPDOp.V_DUAL_MOV_B32,
-               vdstx=v[1], vdsty=v[2], srcx0=s[1], srcy0=s[2], vsrcx1=v[0], vsrcy1=v[0]),
-    ]
-    st = self._run(insts, sgprs={10: 0x40a00000, 11: 0x40e00000, 12: 0x41100000})  # 5.0, 7.0, 9.0
-    self.assertEqual(st._read_vgpr(1, 0), 0x40e00000)  # v[1] = 7.0
-    self.assertEqual(st._read_vgpr(2, 0), 0x41100000)  # v[2] = 9.0
-
-if __name__ == '__main__':
-  unittest.main()

test/mockgpu/amd/amddriver.py

@@ -90,9 +90,9 @@ class AMDDriver(VirtDriver):
   def _prepare_gpu(self, gpu_id):
     self.doorbells[gpu_id] = memoryview(bytearray(0x2000))
     self.gpus[gpu_id] = AMDGPU(gpu_id)
-    # IP versions: rdna3 = GC 11.0.0, NBIF 4.3.0; rdna4 = GC 12.0.0, NBIF 6.3.1
     ip_versions = {"rdna3": {"gc": (11, 0, 0), "sdma": (6, 0, 0), "nbif": (4, 3, 0)},
-                   "rdna4": {"gc": (12, 0, 0), "sdma": (6, 0, 0), "nbif": (6, 3, 1)}}[MOCKGPU_ARCH]
+                   "rdna4": {"gc": (12, 0, 0), "sdma": (6, 0, 0), "nbif": (6, 3, 1)},
+                   "cdna4": {"gc": (9, 5, 0), "sdma": (4, 4, 5), "nbif": (7, 9, 0)}}[MOCKGPU_ARCH]
     def ip_discovery_files(hwid, ver, base_addr):
       p = f'/sys/class/drm/renderD{gpu_id}/device/ip_discovery/die/0/{hwid}/0'
       return [VirtFile(f'/sys/class/drm/renderD{gpu_id}/device/ip_discovery/die/0/{hwid}', functools.partial(DirFileDesc, child_names=['0'])),

test/mockgpu/amd/amdgpu.py

@@ -5,7 +5,7 @@ from tinygrad.helpers import getbits, to_mv, getenv
 from tinygrad.runtime.support import c
 
 MOCKGPU_ARCH = getenv("MOCKGPU_ARCH", "rdna3")
-GFX_TARGET_VERSION = {"rdna3": 110000, "rdna4": 120000}[MOCKGPU_ARCH]
+GFX_TARGET_VERSION = {"rdna3": 110000, "rdna4": 120000, "cdna4": 90500}[MOCKGPU_ARCH]
 import tinygrad.runtime.autogen.amd_gpu as amd_gpu, tinygrad.runtime.autogen.am.pm4_nv as pm4
 
 SDMA_MAX_COPY_SIZE = 0x400000
@@ -106,8 +106,8 @@ class PM4Executor(AMDQueue):
     return (self.rptr[0] - prev_rptr) + executed_in_ib
 
   def _exec_acquire_mem(self, n):
-    assert n == 6
-    for _ in range(7): self._next_dword() # TODO: implement
+    assert n in (5, 6)
+    for _ in range(n + 1): self._next_dword() # TODO: implement
 
   def _exec_release_mem(self, n):
     assert n == 6
@@ -184,6 +184,12 @@ class PM4Executor(AMDQueue):
     args_addr = self.gpu.regs[regCOMPUTE_USER_DATA_0] + (self.gpu.regs[regCOMPUTE_USER_DATA_0 + 1] << 32)
     lc = [self.gpu.regs[i] for i in range(regCOMPUTE_NUM_THREAD_X, regCOMPUTE_NUM_THREAD_X+3)]
     rsrc2 = self.gpu.regs[regCOMPUTE_PGM_RSRC2]
+    # Read all user data registers (hardware loads these directly into s[0:N])
+    user_sgpr_count = (rsrc2 >> 1) & 0x1F  # USER_SGPR_COUNT is bits 1:5
+    user_data = []
+    for i in range(user_sgpr_count):
+      try: user_data.append(self.gpu.regs[regCOMPUTE_USER_DATA_0 + i])
+      except KeyError: user_data.append(0)
 
     prg_sz = 0
     for st,sz in self.gpu.mapped_ranges:
@@ -197,11 +203,12 @@ class PM4Executor(AMDQueue):
     scratch_size = wavesize * 4  # This gives the scratch size per thread (lane)
 
     assert prg_sz > 0, "Invalid prg ptr (not found in mapped ranges)"
-    # Pass valid memory ranges, rsrc2, scratch_size and arch to Python emulator
+    # Pass valid memory ranges, rsrc2, scratch_size, arch, and user data registers to Python emulator
     if hasattr(remu, 'valid_mem_ranges'): remu.valid_mem_ranges = self.gpu.mapped_ranges
     if hasattr(remu, 'rsrc2'): remu.rsrc2 = rsrc2
     if hasattr(remu, 'scratch_size'): remu.scratch_size = scratch_size
     if hasattr(remu, 'arch'): remu.arch = self.gpu.arch
+    if hasattr(remu, 'user_data'): remu.user_data = user_data
     err = remu.run_asm(prg_addr, prg_sz, *gl, *lc, args_addr)
     if err != 0: raise RuntimeError("remu does not support the new instruction introduced in this kernel")
 
@@ -318,7 +325,7 @@ class AMDGPU(VirtGPU):
     self.regs = AMDGPURegisters()
     self.mapped_ranges = set()
     self.queues = []
-    self.arch = MOCKGPU_ARCH
+    self.arch = "cdna" if MOCKGPU_ARCH == "cdna4" else MOCKGPU_ARCH
 
   def map_range(self, vaddr, size): self.mapped_ranges.add((vaddr, size))
   def unmap_range(self, vaddr, size): self.mapped_ranges.remove((vaddr, size))
@@ -329,7 +336,7 @@ class AMDGPU(VirtGPU):
     self.queues.append(SDMAExecutor(self, base, size, rptr, wptr))
     return len(self.queues) - 1
 
-gpu_props = """cpu_cores_count 0
+_gpu_props_rdna = """cpu_cores_count 0
 simd_count 192
 mem_banks_count 1
 caches_count 206
@@ -367,3 +374,44 @@ sdma_fw_version 20
 unique_id 11673270660693242239
 num_xcc 1
 max_engine_clk_ccompute 2400"""
+
+_gpu_props_cdna = """cpu_cores_count 0
+simd_count 304
+mem_banks_count 1
+caches_count 206
+io_links_count 1
+p2p_links_count 5
+cpu_core_id_base 0
+simd_id_base 2147488032
+max_waves_per_simd 16
+lds_size_in_kb 128
+gds_size_in_kb 0
+num_gws 64
+wave_front_size 64
+array_count 16
+simd_arrays_per_engine 4
+cu_per_simd_array 19
+simd_per_cu 2
+max_slots_scratch_cu 32
+gfx_target_version {gfx_target_version}
+vendor_id 4098
+device_id 29772
+location_id 34304
+domain 0
+drm_render_minor {drm_render_minor}
+hive_id 0
+num_sdma_engines 2
+num_sdma_xgmi_engines 0
+num_sdma_queues_per_engine 6
+num_cp_queues 8
+max_engine_clk_fcompute 2100
+local_mem_size 0
+fw_version 2140
+capability 671588992
+debug_prop 1495
+sdma_fw_version 20
+unique_id 11673270660693242239
+num_xcc 1
+max_engine_clk_ccompute 2100"""
+
+gpu_props = _gpu_props_cdna if MOCKGPU_ARCH == "cdna4" else _gpu_props_rdna

test/mockgpu/helpers.py

@@ -21,10 +21,11 @@ class PythonRemu:
   rsrc2: int = 0x19c  # Default: USER_SGPR_COUNT=14, enable X and Y workgroup IDs
   scratch_size: int = 0  # private_segment_fixed_size from kernel descriptor
   arch: str = "rdna3"  # Architecture: rdna3 or rdna4
+  user_data: list[int] = []  # All COMPUTE_USER_DATA registers (loaded into s[0:N])
 
   def run_asm(self, lib: int, lib_sz: int, gx: int, gy: int, gz: int, lx: int, ly: int, lz: int, args_ptr: int) -> int:
     from tinygrad.renderer.amd.emu import run_asm
-    return run_asm(lib, lib_sz, gx, gy, gz, lx, ly, lz, args_ptr, self.rsrc2, self.scratch_size, self.arch)
+    return run_asm(lib, lib_sz, gx, gy, gz, lx, ly, lz, args_ptr, self.rsrc2, self.scratch_size, self.arch, self.user_data)
 
 def _try_dlopen_remu():
   # Use Python emulator only if PYTHON_REMU=1

tinygrad/helpers.py

@@ -189,7 +189,7 @@ CPU_CC, CPU_LLVM, CPU_LVP = ContextVar("CPU_CC", ""), ContextVar("CPU_LLVM", 0),
 NV_CC, NV_PTX, NV_NAK = ContextVar("NV_CC", ""), ContextVar("NV_PTX", 0), ContextVar("NV_NAK", 0)
 CUDA_CC, CUDA_PTX, CUDA_NVCC = ContextVar("CUDA_CC", ""), ContextVar("CUDA_PTX", 0), ContextVar("CUDA_NVCC", 0)
 NULL_IR3, NULL_NAK, NULL_ALLOW_COPYOUT = ContextVar("NULL_IR3", 0), ContextVar("NULL_NAK", 0), ContextVar("NULL_ALLOW_COPYOUT", 0)
-AMD_CC, AMD_LLVM, AMD_HIPCC  = ContextVar("AMD_CC", ""), ContextVar("AMD_LLVM", 0), ContextVar("AMD_HIPCC", 0)
+AMD_CC, AMD_LLVM, AMD_HIPCC, AMD_ISEL, AMD_ASM = ContextVar("AMD_CC", ""), ContextVar("AMD_LLVM", 0), ContextVar("AMD_HIPCC", 0), ContextVar("AMD_ISEL", 0), ContextVar("AMD_ASM", 0)
 QCOM_CC, QCOM_IR3 = ContextVar("QCOM_CC", ""), ContextVar("QCOM_IR3", 0)
 # VIZ implies PROFILE, but you can run PROFILE without VIZ
 VIZ = ContextVar("VIZ", 0)

tinygrad/renderer/amd/dsl.py

@@ -444,6 +444,9 @@ class Inst:
 
   def __eq__(self, other): return type(self) is type(other) and self._raw == other._raw
   def __hash__(self): return hash((type(self), self._raw))
+  def __lt__(self, other):
+    if not isinstance(other, Inst): return NotImplemented
+    return (type(self).__name__, self._raw) < (type(other).__name__, other._raw)
 
   def __repr__(self):
     # collect (repr, is_default) pairs, strip trailing defaults so repr roundtrips with eval

tinygrad/renderer/amd/emu.py

@@ -7,7 +7,7 @@
 #   arg=4: scratch - per-lane scratch memory
 from __future__ import annotations
 import ctypes, functools, re, platform, subprocess, tempfile
-from typing import Any, Callable
+from typing import Callable
 
 # Set/restore DAZ+FTZ (denormals-are-zero + flush-to-zero) to match RDNA3 default float mode
 # x86: MXCSR bits DAZ(6)+FTZ(15), ARM64: FPCR bit FZ(24)
@@ -61,8 +61,10 @@ from tinygrad.engine.realize import get_runner
 from tinygrad.renderer.amd import decode_inst
 from tinygrad.runtime.autogen.amd.rdna3.str_pcode import PCODE as PCODE_RDNA3
 from tinygrad.runtime.autogen.amd.rdna4.str_pcode import PCODE as PCODE_RDNA4
+from tinygrad.runtime.autogen.amd.cdna.str_pcode import PCODE as PCODE_CDNA
 from tinygrad.runtime.autogen.amd.rdna3 import ins as ir3
 from tinygrad.runtime.autogen.amd.rdna4 import ins as ir4
+from tinygrad.runtime.autogen.amd.cdna import ins as irc
 from tinygrad.renderer.amd.dsl import VCC_LO, EXEC_LO, SCC, ttmp
 from tinygrad.runtime.autogen.amd.common import Fmt, OpType
 from tinygrad.renderer.amd.pcode import parse_block, _FUNCS
@@ -160,7 +162,7 @@ _pcode_fixes = {
 
 def _get_pcode_dict(op) -> dict:
   """Return the PCODE dictionary for the given opcode based on its architecture."""
-  return PCODE_RDNA4 if 'rdna4' in type(op).__module__ else PCODE_RDNA3
+  return PCODE_CDNA if 'cdna' in type(op).__module__ else PCODE_RDNA4 if 'rdna4' in type(op).__module__ else PCODE_RDNA3
 
 # Pcode parser
 @functools.cache
@@ -465,8 +467,8 @@ class _Ctx:
     pcode = get_pcode(op)
     vcc_reg = sdst_reg if sdst_reg is not None else VCC_LO.offset
     if 'VCC' not in srcs: srcs['VCC'] = self.rsgpr_dyn(_c(vcc_reg))
-    srcs.update({'EXEC': exec_mask, 'SCC': self.rsgpr_dyn(_c(SCC.offset)), 'laneId': lane,
-                 'ROUND_MODE': _c(0), 'ROUND_TOWARD_ZERO': _c(0)})  # rounding mode: 0=RNE, RTZ constant
+    srcs.update({'EXEC': exec_mask, 'SCC': self.rsgpr_dyn(_c(SCC.offset)), 'laneId': lane, 'VDST': vdst_reg,
+                 'ROUND_MODE': _c(0), 'ROUND_TOWARD_ZERO': _c(0), 'ROUND_NEAREST_EVEN': _c(0)})  # rounding mode constants
     _, assigns = parse_pcode(pcode, srcs)
 
     # For integer ops with clamp, compute overflow using wide arithmetic
@@ -543,10 +545,11 @@ class _Ctx:
 
 def _compile_sopp(inst: ir3.SOPP | ir4.SOPP, ctx: _Ctx) -> UOp:
   simm16 = ctx.inst_field_signed(type(inst).simm16).cast(dtypes.int16)
-  if inst.op in (ir3.SOPPOp.S_ENDPGM, ir4.SOPPOp.S_ENDPGM):
+  if inst.op in (ir3.SOPPOp.S_ENDPGM, ir4.SOPPOp.S_ENDPGM, irc.SOPPOp.S_ENDPGM):
     return UOp.sink(ctx.wsgpr_dyn(_c(PC_LO_IDX), UOp.const(dtypes.uint32, 0xFFFFFFFF)),
                           ctx.wsgpr_dyn(_c(PC_HI_IDX), UOp.const(dtypes.uint32, 0xFFFFFFFF)))
-  if inst.op in (ir3.SOPPOp.S_NOP, ir4.SOPPOp.S_NOP): return UOp.sink(*ctx.inc_pc())  # S_NOP is a no-op
+  # S_NOP and S_WAITCNT are no-ops in emulator (no pipeline/cache to wait on)
+  if inst.op in (ir3.SOPPOp.S_NOP, ir4.SOPPOp.S_NOP, irc.SOPPOp.S_NOP, irc.SOPPOp.S_WAITCNT): return UOp.sink(*ctx.inc_pc())
   # NOTE: we ignore SOPPs without PCODE
   if inst.op in _get_pcode_dict(inst.op):
     pcode = get_pcode(inst.op)
@@ -562,10 +565,7 @@ def _compile_sopp(inst: ir3.SOPP | ir4.SOPP, ctx: _Ctx) -> UOp:
 
 def _compile_smem(inst: ir3.SMEM | ir4.SMEM, ctx: _Ctx) -> UOp:
   # Cache invalidation instructions are no-ops in the emulator (we don't model caches)
-  cache_inv_ops = [ir3.SMEMOp.S_GL1_INV, ir3.SMEMOp.S_DCACHE_INV, ir4.SMEMOp.S_DCACHE_INV]
-  if hasattr(ir4.SMEMOp, 'S_GL1_INV'): cache_inv_ops.append(ir4.SMEMOp.S_GL1_INV)
-  if inst.op in cache_inv_ops:
-    return UOp.sink(*ctx.inc_pc())
+  if '_INV' in inst.op.name: return UOp.sink(*ctx.inc_pc())
   # Dynamic sbase field (bits 5:0) - SGPR pair, field value * 2 = register offset
   sbase = ctx.inst_field(type(inst).sbase) * _c(2)
   # Dynamic sdata field (bits 12:6) - destination SGPR
@@ -573,34 +573,44 @@ def _compile_smem(inst: ir3.SMEM | ir4.SMEM, ctx: _Ctx) -> UOp:
   # RDNA4 uses 'ioffset', RDNA3 uses 'offset' - use type(inst) to get correct field
   offset_field = type(inst).ioffset if hasattr(type(inst), 'ioffset') else type(inst).offset  # type: ignore[union-attr]
   offset = ctx.inst_field_signed(offset_field)  # signed immediate
-  # Dynamic soffset field - SGPR for additional offset (NULL=124 reads as 0)
-  soffset = ctx.inst_field(type(inst).soffset)
-  addr = _u64(ctx.rsgpr_dyn(sbase), ctx.rsgpr_dyn(sbase + _c(1))) + offset.cast(dtypes.uint64) + ctx.rsgpr_dyn(soffset).cast(dtypes.uint64)
+  # Dynamic soffset field - SGPR for additional offset (NULL=124 reads as 0, CDNA soffset_en=0 means no soffset)
+  soffset_val = _c(0).cast(dtypes.uint64)
+  if not (isinstance(inst, irc.SMEM) and not inst.soffset_en):
+    soffset_val = ctx.inst_field(type(inst).soffset)
+    soffset_val = ctx.rsgpr_dyn(soffset_val).cast(dtypes.uint64)
+  addr = _u64(ctx.rsgpr_dyn(sbase), ctx.rsgpr_dyn(sbase + _c(1))) + offset.cast(dtypes.uint64) + soffset_val
   _SMEM_NDWORDS = {ir3.SMEMOp.S_LOAD_B32: 1, ir3.SMEMOp.S_LOAD_B64: 2, ir3.SMEMOp.S_LOAD_B128: 4,
     ir3.SMEMOp.S_LOAD_B256: 8, ir3.SMEMOp.S_LOAD_B512: 16, ir4.SMEMOp.S_LOAD_B32: 1, ir4.SMEMOp.S_LOAD_B64: 2,
-    ir4.SMEMOp.S_LOAD_B96: 3, ir4.SMEMOp.S_LOAD_B128: 4, ir4.SMEMOp.S_LOAD_B256: 8, ir4.SMEMOp.S_LOAD_B512: 16}
+    ir4.SMEMOp.S_LOAD_B96: 3, ir4.SMEMOp.S_LOAD_B128: 4, ir4.SMEMOp.S_LOAD_B256: 8, ir4.SMEMOp.S_LOAD_B512: 16,
+    irc.SMEMOp.S_LOAD_DWORD: 1, irc.SMEMOp.S_LOAD_DWORDX2: 2, irc.SMEMOp.S_LOAD_DWORDX4: 4,
+    irc.SMEMOp.S_LOAD_DWORDX8: 8, irc.SMEMOp.S_LOAD_DWORDX16: 16}
   ndwords = _SMEM_NDWORDS[inst.op]
   stores = [ctx.wsgpr_dyn(sdata_reg + _c(i), ctx.vmem.index((addr + UOp.const(dtypes.uint64, i * 4) >> UOp.const(dtypes.uint64, 2)).cast(dtypes.int)))
             for i in range(ndwords)]
   return UOp.sink(*stores, *ctx.inc_pc())
 
-def _compile_sop(inst: ir3.SOP1 | ir3.SOP2 | ir3.SOPC | ir3.SOPK | ir4.SOP1 | ir4.SOP2 | ir4.SOPC | ir4.SOPK, ctx: _Ctx) -> UOp:
+def _compile_sop(inst: ir3.SOP1|ir3.SOP2|ir3.SOPC|ir3.SOPK|ir4.SOP1|ir4.SOP2|ir4.SOPC|ir4.SOPK|irc.SOP1|irc.SOP2|irc.SOPC|irc.SOPK, ctx: _Ctx) -> UOp:
   bits = inst.canonical_op_bits
   literal = ctx.inst_field(type(inst).literal) if hasattr(type(inst), 'literal') else None  # type: ignore[union-attr]
 
-  if isinstance(inst, (ir3.SOPK, ir4.SOPK)):
+  if isinstance(inst, (ir3.SOPK, ir4.SOPK, irc.SOPK)):
     sdst_off = ctx.inst_field(type(inst).sdst)
     simm16 = ctx.inst_field(type(inst).simm16)
     # Sign-extend simm16
     simm16_sext = simm16.cast(dtypes.int16).cast(dtypes.int32)
-    srcs = {'S0': ctx.rsgpr_dyn(sdst_off), 'SIMM16': simm16_sext, 'D0': ctx.rsgpr_dyn(sdst_off)}
+    # CDNA pcode uses S0 for the immediate in MOVK/MULK/ADDK/CMOVK (where RDNA uses SIMM16),
+    # but S0 = register for CMPK/SETREG. S1 is always the immediate for CDNA CMPK ops.
+    op_name = inst.op.name if hasattr(inst.op, 'name') else ''
+    s0_is_imm = isinstance(inst, irc.SOPK) and 'CMPK' not in op_name and 'SETREG' not in op_name
+    s0_val = simm16_sext if s0_is_imm else ctx.rsgpr_dyn(sdst_off)
+    srcs = {'S0': s0_val, 'SIMM16': simm16_sext, 'S1': simm16_sext, 'D0': ctx.rsgpr_dyn(sdst_off)}
     dst_off, dst_size = sdst_off, 1
-  elif isinstance(inst, (ir3.SOP1, ir4.SOP1)):
+  elif isinstance(inst, (ir3.SOP1, ir4.SOP1, irc.SOP1)):
     sdst_off = ctx.inst_field(type(inst).sdst)
     ssrc0_off = ctx.inst_field(type(inst).ssrc0)
     srcs = {'S0': ctx.rsrc_dyn(ssrc0_off, None, bits['s0'], literal)}
     dst_off, dst_size = sdst_off, bits['d'] // 32
-  elif isinstance(inst, (ir3.SOP2, ir4.SOP2)):
+  elif isinstance(inst, (ir3.SOP2, ir4.SOP2, irc.SOP2)):
     sdst_off = ctx.inst_field(type(inst).sdst)
     ssrc0_off = ctx.inst_field(type(inst).ssrc0)
     ssrc1_off = ctx.inst_field(type(inst).ssrc1)
@@ -608,7 +618,7 @@ def _compile_sop(inst: ir3.SOP1 | ir3.SOP2 | ir3.SOPC | ir3.SOPK | ir4.SOP1 | ir
             'S1': ctx.rsrc_dyn(ssrc1_off, None, bits['s1'], literal)}
     if literal is not None: srcs['SIMM32'] = literal
     dst_off, dst_size = sdst_off, bits['d'] // 32
-  elif isinstance(inst, (ir3.SOPC, ir4.SOPC)):
+  elif isinstance(inst, (ir3.SOPC, ir4.SOPC, irc.SOPC)):
     ssrc0_off = ctx.inst_field(type(inst).ssrc0)
     ssrc1_off = ctx.inst_field(type(inst).ssrc1)
     srcs = {'S0': ctx.rsrc_dyn(ssrc0_off, None, bits['s0'], literal),
@@ -619,7 +629,7 @@ def _compile_sop(inst: ir3.SOP1 | ir3.SOP2 | ir3.SOPC | ir3.SOPK | ir4.SOP1 | ir
 
   return ctx.compile_sop_pcode(inst.op, srcs, dst_off, dst_size)
 
-def _compile_vop12(inst: ir3.VOP1 | ir3.VOP1_SDST | ir3.VOP2 | ir4.VOP1 | ir4.VOP1_SDST | ir4.VOP2, ctx: _Ctx) -> UOp:
+def _compile_vop12(inst: ir3.VOP1 | ir3.VOP1_SDST | ir3.VOP2 | ir4.VOP1 | ir4.VOP1_SDST | ir4.VOP2 | irc.VOP1 | irc.VOP2, ctx: _Ctx) -> UOp:
   op_name = _op_name(inst)
   if op_name in ('V_READFIRSTLANE_B32_E32', 'V_PERMLANE64_B32_E32'): return ctx.compile_lane_pcode(inst.op, inst)
   lane, exec_mask, bits = ctx.range(), ctx.rsgpr_dyn(_c(EXEC_LO.offset)), inst.canonical_op_bits
@@ -628,7 +638,7 @@ def _compile_vop12(inst: ir3.VOP1 | ir3.VOP1_SDST | ir3.VOP2 | ir4.VOP1 | ir4.VO
   write_hi_half = bits['d'] == 16 and (vdst_reg >= _c(128))
   if isinstance(write_hi_half, UOp): vdst_reg = write_hi_half.where(vdst_reg - _c(128), vdst_reg)
   elif write_hi_half: vdst_reg -= 128
-  if isinstance(inst, (ir3.VOP1, ir4.VOP1)):
+  if isinstance(inst, (ir3.VOP1, ir4.VOP1, irc.VOP1)):
     # Handle VOP1 hi-half source operand (src0 >= v[128] for 16-bit ops)
     src0_off = ctx.inst_field(type(inst).src0)
     s0 = ctx.rsrc_dyn(src0_off, lane, bits['s0'], literal)
@@ -654,12 +664,13 @@ def _compile_vop12(inst: ir3.VOP1 | ir3.VOP1_SDST | ir3.VOP2 | ir4.VOP1 | ir4.VO
       s0 = src0_hi.where(_hi16(ctx.rvgpr_dyn(src0_reg, lane)), s0)
     srcs = {'S0': s0, 'S1': s1, 'D0': d0}
     if inst.op in (ir3.VOP2Op.V_FMAAK_F32_E32, ir3.VOP2Op.V_FMAMK_F32_E32, ir3.VOP2Op.V_FMAAK_F16_E32,
-                   ir3.VOP2Op.V_FMAMK_F16_E32):
+                   ir3.VOP2Op.V_FMAMK_F16_E32, irc.VOP2Op.V_FMAAK_F32_E32, irc.VOP2Op.V_FMAMK_F32_E32):
       assert literal is not None
       srcs['SIMM32'] = literal
   return ctx.compile_vop_pcode(inst.op, srcs, lane, vdst_reg, exec_mask, opsel_dst_hi=write_hi_half)
 
-def _compile_vopc(inst: ir3.VOPC | ir3.VOP3 | ir4.VOPC | ir4.VOP3, ctx: _Ctx, opsel: int = 0, abs_bits: int = 0, neg_bits: int = 0) -> UOp:
+def _compile_vopc(inst: ir3.VOPC|ir3.VOP3|ir4.VOPC|ir4.VOP3|irc.VOPC|irc.VOP3, ctx: _Ctx,
+                  opsel: int = 0, abs_bits: int = 0, neg_bits: int = 0) -> UOp:
   exec_mask, op_name, bits = ctx.rsgpr_dyn(_c(EXEC_LO.offset)), _op_name(inst), inst.canonical_op_bits
   is_cmpx, is_vopc = 'CMPX' in op_name, hasattr(inst, 'vsrc1')  # is_vopc: e32 vs e64
 
@@ -707,7 +718,7 @@ def _compile_vopc(inst: ir3.VOPC | ir3.VOP3 | ir4.VOPC | ir4.VOP3, ctx: _Ctx, op
     stores = [ctx.wsgpr_dyn(dst_off, new_result)] if not is_vopc else [ctx.wsgpr_dyn(_c(VCC_LO.offset), new_result)]
   return UOp.sink(*stores, *ctx.inc_pc())
 
-def _compile_vop3(inst: ir3.VOP3 | ir4.VOP3, ctx: _Ctx) -> UOp:
+def _compile_vop3(inst: ir3.VOP3 | ir4.VOP3 | irc.VOP3, ctx: _Ctx) -> UOp:
   exec_mask = ctx.rsgpr_dyn(_c(EXEC_LO.offset))
   bits = inst.canonical_op_bits
   opsel, op_name = getattr(inst, 'opsel', 0) or 0, _op_name(inst)
@@ -741,13 +752,13 @@ def _compile_vop3(inst: ir3.VOP3 | ir4.VOP3, ctx: _Ctx) -> UOp:
   src1 = _apply_src_mods(src1, 1, abs_bits, neg_bits, bits['s1'])
   src2 = _apply_src_mods(src2, 2, abs_bits, neg_bits, bits['s2'])
   srcs = {'S0': src0, 'S1': src1, 'S2': src2}
-  if inst.op in (ir3.VOP3Op.V_CNDMASK_B32_E64, ir3.VOP3Op.V_CNDMASK_B16) and src2 is not None: srcs['VCC'] = src2
+  if inst.op in (ir3.VOP3Op.V_CNDMASK_B32_E64, ir3.VOP3Op.V_CNDMASK_B16, irc.VOP3Op.V_CNDMASK_B32_E64) and src2 is not None: srcs['VCC'] = src2
   # FMAC instructions need D0 (accumulator) from destination register
   if 'FMAC' in op_name: srcs['D0'] = ctx.rvgpr_dyn(vdst_reg, lane)
   opsel_dst_hi = bool(opsel & 0b1000) and bits['d'] == 16
   return ctx.compile_vop_pcode(inst.op, srcs, lane, vdst_reg, exec_mask, opsel_dst_hi=opsel_dst_hi, clmp=getattr(inst, 'clmp', 0))
 
-def _compile_vop3sd(inst: ir3.VOP3SD | ir4.VOP3SD, ctx: _Ctx) -> UOp:
+def _compile_vop3sd(inst: ir3.VOP3SD | ir4.VOP3SD | irc.VOP3SD, ctx: _Ctx) -> UOp:
   exec_mask = ctx.rsgpr_dyn(_c(EXEC_LO.offset))
   bits, pcode, ops = inst.canonical_op_bits, get_pcode(inst.op), inst.canonical_operands
 
@@ -806,7 +817,7 @@ def _compile_vop3sd(inst: ir3.VOP3SD | ir4.VOP3SD, ctx: _Ctx) -> UOp:
   else:
     return ctx.compile_vop_pcode(inst.op, srcs, lane, vdst_reg, exec_mask, sdst_reg=inst.sdst.offset)
 
-def _compile_wmma(inst: ir3.VOP3P | ir4.VOP3P, ctx: _Ctx) -> UOp:
+def _compile_wmma(inst: ir3.VOP3P | ir4.VOP3P | irc.VOP3P, ctx: _Ctx) -> UOp:
   op_name = _op_name(inst)
   exec_mask = ctx.rsgpr_dyn(_c(EXEC_LO.offset))
   vdst_reg = ctx.inst_field(type(inst).vdst)
@@ -839,14 +850,15 @@ def _compile_wmma(inst: ir3.VOP3P | ir4.VOP3P, ctx: _Ctx) -> UOp:
     stores = [ctx.wvgpr_dyn(vdst_reg + _c(i // 32), UOp.const(dtypes.int, i % 32), mat_d[i].bitcast(dtypes.uint32), exec_mask) for i in range(256)]
   return UOp.sink(*stores, *ctx.inc_pc())
 
-def _compile_vop3p(inst: ir3.VOP3P | ir4.VOP3P, ctx: _Ctx) -> UOp:
+def _compile_vop3p(inst: ir3.VOP3P | ir4.VOP3P | irc.VOP3P, ctx: _Ctx) -> UOp:
   op_name = _op_name(inst)
   if 'WMMA' in op_name and ('16X16X16_F16' in op_name or '16X16X16_BF16' in op_name): return _compile_wmma(inst, ctx)
 
   lane = ctx.range()
   exec_mask = ctx.rsgpr_dyn(_c(EXEC_LO.offset))
   vdst_reg = ctx.inst_field(type(inst).vdst)
-  do_cast = any(x in op_name for x in ('F16', 'F32', 'BF16')) and 'IU' not in op_name
+  is_pk_f32 = 'PK' in op_name and 'F32' in op_name and 'MOV' not in op_name  # CDNA packed F32 ops
+  do_cast = any(x in op_name for x in ('F16', 'F32', 'BF16')) and 'IU' not in op_name and not is_pk_f32
   src0 = ctx.rsrc_dyn(ctx.inst_field(type(inst).src0), lane, 16, do_cast=do_cast)
   src1 = ctx.rsrc_dyn(ctx.inst_field(type(inst).src1), lane, 16, do_cast=do_cast)
   src2 = ctx.rsrc_dyn(ctx.inst_field(type(inst).src2), lane, 16, do_cast=do_cast)
@@ -854,7 +866,30 @@ def _compile_vop3p(inst: ir3.VOP3P | ir4.VOP3P, ctx: _Ctx) -> UOp:
   opsel_hi2 = getattr(inst, 'opsel_hi2', 1) if getattr(inst, 'opsel_hi2', 1) is not None else 1
   neg, neg_hi = getattr(inst, 'neg', 0) or 0, getattr(inst, 'neg_hi', 0) or 0
 
-  if 'FMA_MIX' in op_name:
+  if is_pk_f32:
+    # CDNA packed F32: read 32-bit sources, build 64-bit packed values using opsel.
+    # For VGPRs: opsel selects between v[reg] (0) and v[reg+1] (1) for each half.
+    # For SGPR pairs (off < 128): s[N] = lo float32, s[N+1] = hi float32.
+    # For inline constants (128 <= off < 256): broadcast same value to both halves.
+    src_offs = [ctx.inst_field(type(inst).src0), ctx.inst_field(type(inst).src1), ctx.inst_field(type(inst).src2)]
+    def build_pk_f32(src_lo: UOp, src_off: UOp, opsel_lo: int, opsel_hi_bit: int, neg_lo: int, neg_hi_bit: int) -> UOp:
+      is_vgpr = src_off >= _c(256)
+      vgpr_lo = ctx.rvgpr_dyn(src_off - _c(256), lane) if lane is not None else _c(0)
+      vgpr_hi = ctx.rvgpr_dyn(src_off - _c(256) + _c(1), lane) if lane is not None else _c(0)
+      # For SGPR pairs, opsel selects between s[N] (0) and s[N+1] (1); inline constants always broadcast.
+      is_sgpr_pair = src_off < _c(128)
+      sgpr_hi = ctx.rsgpr_dyn(src_off + _c(1), is_sgpr_pair)
+      scalar_lo_sel = src_lo if not opsel_lo else is_sgpr_pair.where(sgpr_hi, src_lo)
+      scalar_hi_sel = src_lo if not opsel_hi_bit else is_sgpr_pair.where(sgpr_hi, src_lo)
+      lo = is_vgpr.where(vgpr_hi if opsel_lo else vgpr_lo, scalar_lo_sel)
+      hi = is_vgpr.where(vgpr_hi if opsel_hi_bit else vgpr_lo, scalar_hi_sel)
+      if neg_lo: lo = lo ^ UOp.const(dtypes.uint32, 0x80000000)
+      if neg_hi_bit: hi = hi ^ UOp.const(dtypes.uint32, 0x80000000)
+      return _u64(lo, hi)
+    srcs = {'S0': build_pk_f32(src0, src_offs[0], opsel & 1, opsel_hi & 1, neg & 1, neg_hi & 1),
+            'S1': build_pk_f32(src1, src_offs[1], opsel & 2, opsel_hi & 2, neg & 2, neg_hi & 2),
+            'S2': build_pk_f32(src2, src_offs[2], opsel & 4, 1 if opsel_hi2 else 0, neg & 4, neg_hi & 4)}
+  elif 'FMA_MIX' in op_name:
     combined_opsel_hi = (opsel_hi & 0x3) | ((opsel_hi2 & 0x1) << 2)
     # For FMA_MIX: neg_hi is ABS (not neg!), neg is actual negation
     def apply_abs(v, bit, opsel_hi_bit, opsel_bit):
@@ -924,13 +959,18 @@ def _compile_vopd(inst: ir3.VOPD | ir4.VOPD, ctx: _Ctx) -> UOp:
       if dest.startswith('D0'): all_stores.append(ctx.wvgpr_dyn(vdst_reg, lane, _val_to_u32(val), exec_mask, after=srcy1))
   return UOp.sink(UOp.group(*all_stores).end(lane), *ctx.inc_pc())
 
-def _compile_mem_op(inst: ir3.DS | ir3.FLAT | ir3.GLOBAL | ir3.SCRATCH | ir4.DS | ir4.VFLAT | ir4.VGLOBAL | ir4.VSCRATCH, ctx: _Ctx) -> UOp:
+def _compile_mem_op(inst: ir3.DS|ir3.FLAT|ir3.GLOBAL|ir3.SCRATCH|ir4.DS|ir4.VFLAT|ir4.VGLOBAL|ir4.VSCRATCH
+                    |irc.DS|irc.FLAT|irc.GLOBAL|irc.SCRATCH, ctx: _Ctx) -> UOp:
   """Unified memory operation compiler for DS, FLAT, GLOBAL, SCRATCH."""
   exec_mask, op_name = ctx.rsgpr_dyn(_c(EXEC_LO.offset)), _op_name(inst)
   pcode = get_pcode(inst.op)
+  # CDNA pcode uses CalcGlobalAddr/CalcDsAddr to compute address from raw components, but make_addr already handles this.
+  # Strip the addr computation line and use pre-computed ADDR directly (rename 'addr' -> 'ADDR' in remaining pcode).
+  if isinstance(inst, (irc.GLOBAL, irc.FLAT, irc.SCRATCH, irc.DS)) and 'Calc' in pcode and 'Addr' in pcode:
+    pcode = re.sub(r'addr\s*=\s*Calc\w+Addr\([^)]*\)\s*;?\n?', '', pcode).replace('MEM[addr', 'MEM[ADDR')
 
-  is_lds = isinstance(inst, (ir3.DS, ir4.DS))
-  is_scratch = isinstance(inst, (ir3.SCRATCH, ir4.VSCRATCH))
+  is_lds = isinstance(inst, (ir3.DS, ir4.DS, irc.DS))
+  is_scratch = isinstance(inst, (ir3.SCRATCH, ir4.VSCRATCH, irc.SCRATCH))
   mem = ctx.lds if is_lds else ctx.scratch if is_scratch else ctx.vmem
   addr_shift = UOp.const(dtypes.uint32 if is_lds else dtypes.uint64, 2)
 
@@ -1038,7 +1078,7 @@ def _compile_mem_op(inst: ir3.DS | ir3.FLAT | ir3.GLOBAL | ir3.SCRATCH | ir4.DS
     if 'STORE' in op_name and data_bits_mem >= 64:
       vdata = vdata | (ctx.rvgpr_dyn(vdata_reg + _c(1), lane).cast(dtypes.uint64) << UOp.const(dtypes.uint64, 32))
     srcs = {'ADDR': addr, 'VDATA': vdata, '_vmem': mem, '_active': active,
-            'laneId': lane, 'v_addr': vaddr_base, 's_saddr': saddr_base}
+            'laneId': lane, 'v_addr': vaddr_base, 's_saddr': saddr_base, 'SADDR': saddr_base, 'OFFSET': offset}
     for i in range(data_bits_mem // 32):
       srcs[f'VDATA{i}'] = ctx.rvgpr_dyn(vdata_reg + _c(i), lane) if 'STORE' in op_name else UOp.const(dtypes.uint32, 0)
     return srcs
@@ -1075,7 +1115,7 @@ def _compile_mem_op(inst: ir3.DS | ir3.FLAT | ir3.GLOBAL | ir3.SCRATCH | ir4.DS
       return UOp.sink(*ended, *ctx.inc_pc())
 
   # Standard path: single lane range
-  writes_return_data = '_RTN' in op_name or (is_lds and op_name.startswith('DS_LOAD')) or bool(is_atomic and glc)
+  writes_return_data = '_RTN' in op_name or (is_lds and (op_name.startswith('DS_LOAD') or op_name.startswith('DS_READ'))) or bool(is_atomic and glc)
   lane = ctx.range()
   active = _lane_active(exec_mask, lane)
   pcode_vars, assigns = parse_pcode(pcode, make_srcs(lane))
@@ -1099,6 +1139,11 @@ _INST_HANDLERS: dict[type, Callable[..., UOp]] = {
   ir4.VOP1: _compile_vop12, ir4.VOP1_SDST: _compile_vop12, ir4.VOP2: _compile_vop12, ir4.VOPC: _compile_vopc, ir4.VOP3: _compile_vop3,
   ir4.VOP3_SDST: _compile_vop3, ir4.VOP3SD: _compile_vop3sd, ir4.VOP3P: _compile_vop3p, ir4.VOPD: _compile_vopd,
   ir4.DS: _compile_mem_op, ir4.VFLAT: _compile_mem_op, ir4.VGLOBAL: _compile_mem_op, ir4.VSCRATCH: _compile_mem_op,
+  # CDNA instruction classes
+  irc.SOPP: _compile_sopp, irc.SMEM: _compile_smem, irc.SOP1: _compile_sop, irc.SOP2: _compile_sop, irc.SOPC: _compile_sop, irc.SOPK: _compile_sop,
+  irc.VOP1: _compile_vop12, irc.VOP2: _compile_vop12, irc.VOPC: _compile_vopc, irc.VOP3: _compile_vop3,
+  irc.VOP3_SDST: _compile_vop3, irc.VOP3SD: _compile_vop3sd, irc.VOP3P: _compile_vop3p,
+  irc.DS: _compile_mem_op, irc.FLAT: _compile_mem_op, irc.GLOBAL: _compile_mem_op, irc.SCRATCH: _compile_mem_op,
 }
 
 # ═══════════════════════════════════════════════════════════════════════════════
@@ -1116,7 +1161,7 @@ def _get_runner(inst_bytes: bytes, arch: str = "rdna3"):
 
   # Check if instruction matches any cached canonical pattern
   for base, mask, size, runner in _canonical_runner_cache:
-    if inst_size == size and (inst_int & mask) == base: return runner, False
+    if inst_size == size and (inst_int & mask) == base: return runner
 
   # Look up handler by type, falling back to base classes for _LIT variants
   handler = _INST_HANDLERS.get(type(inst))
@@ -1136,30 +1181,17 @@ def _get_runner(inst_bytes: bytes, arch: str = "rdna3"):
   with Context(NOOPT=1, CHECK_OOB=0, TUPLE_ORDER=0, EMULATED_DTYPES=""):
     runner = get_runner('CPU', sink)
   _canonical_runner_cache.append((base, mask, size, runner))
-  return runner, True
+  return runner
 
-@functools.cache
-def decode_program(data: bytes, arch: str = "rdna3") -> dict[int, tuple[str, Callable, list[int], Any]]:
-  """Decode program to {pc: (name, fxn, globals, runner)}."""
-  result: dict[int, tuple[str, Callable, list[int], Any]] = {}
-  i = 0
-  while i < len(data):
-    inst = decode_inst(data[i:], arch)
-    if hasattr(inst, 'op') and inst.op in (ir3.SOPPOp.S_CODE_END, ir4.SOPPOp.S_CODE_END): break
-    try:
-      runner, is_new = _get_runner(bytes(data[i:i + inst.size() + 4]), arch)
-      if DEBUG >= 3:
-        try: inst_str = repr(inst)
-        except Exception: inst_str = f"<{type(inst).__name__} at PC={i}>"
-        msg = f"[emu] PC={i}: {inst_str}"
-        print(colored(msg, 'green') if is_new else msg)
-      result[i] = (runner.p.function_name, runner._prg.fxn, runner.p.globals, runner)
-    except Exception as e:
-      try: inst_str = repr(inst)
-      except Exception: inst_str = f"<{type(inst).__name__}>"
-      raise RuntimeError(f"[emu] Failed to compile PC={i} {inst_str}: {type(e).__name__}: {e}") from e
-    i += inst.size()
-  return result
+def _decode_at(pc: int, arch: str):
+  """Decode and compile instruction at absolute address pc. Returns CompiledRunner."""
+  inst_bytes = bytes((ctypes.c_char * 16).from_address(pc).raw)
+  inst = decode_inst(inst_bytes, arch)
+  try: return _get_runner(bytes(inst_bytes[:inst.size() + 4]), arch)
+  except Exception as e:
+    try: inst_str = repr(inst)
+    except Exception: inst_str = f"<{type(inst).__name__}>"
+    raise RuntimeError(f"[emu] Failed to compile {inst_str}: {type(e).__name__}: {e}") from e
 
 # ═══════════════════════════════════════════════════════════════════════════════
 # WAVE STATE
@@ -1206,10 +1238,9 @@ class WaveState:
 # ═══════════════════════════════════════════════════════════════════════════════
 
 def run_asm(lib: int, lib_sz: int, gx: int, gy: int, gz: int, lx: int, ly: int, lz: int, args_ptr: int, rsrc2: int = 0x19c,
-            scratch_size: int = 0, arch: str = "rdna3") -> int:
+            scratch_size: int = 0, arch: str = "rdna3", user_data: list[int]|None = None) -> int:
   """Execute AMD assembly program. scratch_size is private_segment_fixed_size from kernel descriptor (per-lane)."""
-  program_raw = decode_program(bytes((ctypes.c_char * lib_sz).from_address(lib).raw), arch)
-  program = {lib + offset: val for offset, val in program_raw.items()}  # Remap to actual addresses
+  program: dict[int, tuple[Callable, list[int]]] = {}  # lazily populated: pc -> (fxn, globals) extracted from runner
   lds_size = ((rsrc2 & hsa.AMD_COMPUTE_PGM_RSRC_TWO_GRANULATED_LDS_SIZE) >> hsa.AMD_COMPUTE_PGM_RSRC_TWO_GRANULATED_LDS_SIZE_SHIFT) * 512
   total_threads = lx * ly * lz
 
@@ -1226,8 +1257,12 @@ def run_asm(lib: int, lib_sz: int, gx: int, gy: int, gz: int, lx: int, ly: int,
           for wave_start in range(0, total_threads, WAVE_SIZE):
             n_lanes, st = min(WAVE_SIZE, total_threads - wave_start), WaveState(min(WAVE_SIZE, total_threads - wave_start))
             st.pc = lib  # Set PC to code base address
-            st._write_sgpr(0, args_ptr & MASK32)
-            st._write_sgpr(1, (args_ptr >> 32) & MASK32)
+            # Initialize user SGPRs: hardware loads COMPUTE_USER_DATA registers directly into s[0:N]
+            if user_data:
+              for i, val in enumerate(user_data): st._write_sgpr(i, val)
+            else:
+              st._write_sgpr(0, args_ptr & MASK32)
+              st._write_sgpr(1, (args_ptr >> 32) & MASK32)
 
             # Workgroup IDs in SGPRs after user SGPRs
             sgpr_idx = (rsrc2 & hsa.AMD_COMPUTE_PGM_RSRC_TWO_USER_SGPR_COUNT) >> hsa.AMD_COMPUTE_PGM_RSRC_TWO_USER_SGPR_COUNT_SHIFT
@@ -1255,13 +1290,16 @@ def run_asm(lib: int, lib_sz: int, gx: int, gy: int, gz: int, lx: int, ly: int,
                       ctypes.c_uint64(vmem_buf._buf.va_addr), ctypes.c_uint64(lds_buf._buf.va_addr),
                       ctypes.c_uint64(scratch_buf._buf.va_addr if scratch_buf else 0)]
             for inst_count in range(1_000_000):
-              if (pc := st.pc) == 0xFFFFFFFFFFFFFFFF or pc not in program: break
-              name, fxn, globals_list, _ = program[pc]
-              assert fxn is not None, f"[emu] No fxn for {name} at PC={pc}"
-              assert 4 not in globals_list or scratch_buf, f"SCRATCH instruction {name} but scratch_size=0"
-              if DEBUG >= 6:
-                inst = decode_inst(bytes((ctypes.c_char * 12).from_address(pc).raw), arch)
-                print(f"[emu] exec PC={pc:X}: {inst!r}")
+              if (pc := st.pc) == 0xFFFFFFFFFFFFFFFF: break
+              if pc not in program:
+                prev_len = len(_canonical_runner_cache)
+                runner = _decode_at(pc, arch)
+                program[pc] = (runner._prg.fxn, runner.p.globals)
+                if DEBUG >= 3:
+                  inst = decode_inst(bytes((ctypes.c_char * 16).from_address(pc).raw), arch)
+                  msg = f"[emu] PC={pc - lib}: {inst!r}"
+                  print(colored(msg, 'green') if len(_canonical_runner_cache) > prev_len else msg)
+              fxn, globals_list = program[pc]
               fxn(*[c_bufs[g] for g in globals_list])
             else: raise RuntimeError("exceeded 1M instructions, likely infinite loop")
   return 0

tinygrad/renderer/amd/isel.py

@@ -0,0 +1,325 @@
+# Instruction selection for AMD GPUs via pcode-derived PatternMatcher
+# Parses AMD pcode specs into UOp templates, normalizes them, and converts to UPat patterns
+# that rewrite renderer-level UOps into Ops.INS with arg=Inst objects
+
+import functools
+from tinygrad.uop.ops import UOp, UPat, Ops, PatternMatcher, GroupOp
+from tinygrad.dtype import dtypes, DType
+from extra.assembly.amd.emu import parse_pcode
+from extra.assembly.amd.autogen.rdna3.str_pcode import PCODE
+from extra.assembly.amd.autogen.rdna3 import ins as rdna3_ins
+
+# sentinel UOps representing source operands (typed as u32, like real registers)
+_SENTINEL = {f'S{i}': UOp(Ops.DEFINE_VAR, dtypes.uint32, arg=(f'S{i}', 0, 0xFFFFFFFF)) for i in range(4)}
+_SENTINEL_SET = set(_SENTINEL.values())
+
+# only parse ALU-relevant opcode types (memory ops need different sentinels)
+_ALU_ENUM_TYPES = frozenset({'SOP1Op', 'SOP2Op', 'SOPCOp', 'SOPKOp', 'VOP1Op', 'VOP2Op',
+                             'VOP3Op', 'VOP3POp', 'VOP3SDOp', 'VOPCOp', 'VINTERPOp'})
+# SOP types use SGPRs — skip for LLVM inline asm renderer (VGPR-only)
+_SOP_ENUM_TYPES = frozenset({'SOP1Op', 'SOP2Op', 'SOPCOp', 'SOPKOp'})
+
+# opcode enum class name -> list of Inst class suffixes to try
+_VARIANT_SUFFIXES = ['', '_SDST']
+
+def make_inst(opcode):
+  """Create an Inst object with just the opcode set (registers defaulted)."""
+  base_name = type(opcode).__name__[:-2]
+  for suffix in _VARIANT_SUFFIXES:
+    cls = getattr(rdna3_ins, base_name + suffix, None)
+    if cls is None: continue
+    try: return cls(op=opcode)
+    except (RuntimeError, TypeError): continue
+  raise RuntimeError(f"no Inst class found for {opcode}")
+
+# ═══════════════════════════════════════════════════════════════
+# Normalization: strip register-model artifacts from pcode UOps
+# ═══════════════════════════════════════════════════════════════
+
+def normalize(uop, _cache=None):
+  """Strip register-model artifacts from pcode UOps to match renderer-level UOps.
+
+  Pcode models registers as typeless u32 words and uses BITCAST/CAST to reinterpret.
+  The renderer's UOps are natively typed — we strip these artifacts:
+    - BITCAST(f32, sentinel_u32) -> sentinel typed as f32
+    - CAST(i32, sentinel_u32) -> sentinel typed as i32 (same-size reinterpret)
+    - CAST(u64, sentinel_u32) -> sentinel typed as u64 (widening for 64-bit ops)
+    - BITCAST(T, x) where x.dtype == T -> x (identity bitcast)
+    - AND(x, mask) where mask is shift masking -> x (hardware does this implicitly)
+  """
+  if _cache is None: _cache = {}
+  if id(uop) in _cache: return _cache[id(uop)]
+
+  # first recurse so children are normalized before we check patterns
+  new_src = tuple(normalize(s, _cache) for s in uop.src)
+  uop = uop if new_src == uop.src else uop.replace(src=new_src)
+
+  # BITCAST or CAST on a sentinel -> sentinel with target dtype
+  if uop.op in (Ops.BITCAST, Ops.CAST) and len(uop.src) == 1 and uop.src[0] in _SENTINEL_SET:
+    result = uop.src[0].replace(dtype=uop.dtype)
+    _cache[id(uop)] = result
+    return result
+
+  # identity BITCAST: BITCAST(T, x) where x already has dtype T
+  if uop.op == Ops.BITCAST and len(uop.src) == 1 and uop.src[0].dtype == uop.dtype:
+    _cache[id(uop)] = uop.src[0]
+    return uop.src[0]
+
+  # shift masking: AND(sentinel, 31) or AND(sentinel, 63) -> sentinel (hardware masks shift amounts)
+  if uop.op == Ops.AND and len(uop.src) == 2:
+    if uop.src[1].op == Ops.CONST and uop.src[1].arg in (31, 63) and uop.src[0].op == Ops.DEFINE_VAR:
+      _cache[id(uop)] = uop.src[0]
+      return uop.src[0]
+
+  _cache[id(uop)] = uop
+  return uop
+
+def _count_nodes(uop, _seen=None):
+  if _seen is None: _seen = set()
+  if id(uop) in _seen: return 0
+  _seen.add(id(uop))
+  return 1 + sum(_count_nodes(s, _seen) for s in uop.src)
+
+# ═══════════════════════════════════════════════════════════════
+# UOp template -> UPat conversion
+# ═══════════════════════════════════════════════════════════════
+
+def uop_to_upat(uop, _seen=None):
+  """Convert a normalized UOp template into a matchable UPat pattern."""
+  if _seen is None: _seen = {}
+  if id(uop) in _seen: return _seen[id(uop)]
+  if uop.op == Ops.DEFINE_VAR and isinstance(uop.arg, tuple) and uop.arg[0] in _SENTINEL:
+    result = UPat.var(uop.arg[0], dtype=uop.dtype)
+    _seen[id(uop)] = result
+    return result
+  if uop.op in (Ops.CONST, Ops.VCONST):
+    result = UPat(uop.op, uop.dtype, arg=uop.arg)
+    _seen[id(uop)] = result
+    return result
+  src = tuple(uop_to_upat(s, _seen) for s in uop.src) if uop.src else None
+  result = UPat(uop.op, uop.dtype, src=src)
+  _seen[id(uop)] = result
+  return result
+
+# ═══════════════════════════════════════════════════════════════
+# Pattern classification and selection
+# ═══════════════════════════════════════════════════════════════
+
+def _select_best_opcode(opcodes):
+  """Prefer shorter encodings: VOP1/VOP2 > SOP > VOP3/VOPC."""
+  _PREF = {'VOP1Op': 0, 'VOP2Op': 0, 'SOP1Op': 1, 'SOP2Op': 1, 'SOPCOp': 1, 'VOPCOp': 2, 'VOP3Op': 3, 'VOP3SDOp': 3, 'VOP3POp': 4}
+  return min(opcodes, key=lambda oc: (_PREF.get(type(oc).__name__, 9), oc.value))
+
+def _is_direct_alu(norm_uop):
+  """Check if normalized UOp is a direct ALU: op(sentinels...) with no intermediate ops."""
+  if norm_uop.op not in GroupOp.ALU and norm_uop.op not in {Ops.CAST, Ops.BITCAST}: return False
+  return all(s.op == Ops.DEFINE_VAR for s in norm_uop.src)
+
+def _pattern_key(uop, _seen=None):
+  """Structural fingerprint for a normalized UOp template (sentinels become var placeholders)."""
+  if _seen is None: _seen = {}
+  if id(uop) in _seen: return _seen[id(uop)]
+  if uop.op == Ops.DEFINE_VAR and isinstance(uop.arg, tuple) and uop.arg[0] in _SENTINEL:
+    result = f'var({uop.arg[0]},{uop.dtype})'
+  elif uop.op in (Ops.CONST, Ops.VCONST):
+    result = f'const({uop.arg},{uop.dtype})'
+  else:
+    children = ','.join(_pattern_key(s, _seen) for s in uop.src)
+    result = f'{uop.op}({uop.dtype},{children})'
+  _seen[id(uop)] = result
+  return result
+
+def _runtime_key(uop, _var_counter=None, _seen=None):
+  """Compute a structural key from a matched UOp at runtime (real data, not sentinels).
+  Leaf UOps (non-ALU with no recognized children) are treated as variables."""
+  if _seen is None: _seen = {}
+  if _var_counter is None: _var_counter = [0]
+  uid = id(uop)
+  if uid in _seen: return _seen[uid]
+  _ALU_OPS = GroupOp.ALU | {Ops.CAST, Ops.BITCAST, Ops.WHERE}
+  if uop.op in (Ops.CONST, Ops.VCONST):
+    result = f'const({uop.arg},{uop.dtype})'
+  elif uop.op not in _ALU_OPS:
+    result = f'var(S{_var_counter[0]},{uop.dtype})'
+    _var_counter[0] += 1
+  else:
+    children = ','.join(_runtime_key(s, _var_counter, _seen) for s in uop.src)
+    result = f'{uop.op}({uop.dtype},{children})'
+  _seen[uid] = result
+  return result
+
+# ═══════════════════════════════════════════════════════════════
+# Build tables from pcode
+# ═══════════════════════════════════════════════════════════════
+
+def _parse_pcode_patterns(pcode_dict, vgpr_only=False):
+  """Parse ALU pcode entries, normalize, and categorize into direct vs structural."""
+  # key: (op, src_dtypes_tuple, dst_dtype) -> [opcodes]
+  direct: dict[tuple, list] = {}
+  structural: list[tuple] = []  # [(opcode, norm_uop)]
+  allowed_types = _ALU_ENUM_TYPES - _SOP_ENUM_TYPES if vgpr_only else _ALU_ENUM_TYPES
+
+  for opcode, pcode_str in pcode_dict.items():
+    if type(opcode).__name__ not in allowed_types: continue
+    try: env, assigns = parse_pcode(pcode_str, dict(_SENTINEL))
+    except Exception: continue
+    d0 = next(((n, u) for n, u in assigns if n.startswith('D0')), None)
+    if d0 is None: continue
+    _, uop = d0
+    if _count_nodes(uop) > 5: continue
+    norm = normalize(uop)
+    if norm.op == Ops.DEFINE_VAR or norm.op in (Ops.CONST, Ops.VCONST): continue
+
+    if _is_direct_alu(norm):
+      src_dtypes = tuple(s.dtype for s in norm.src)
+      key = (norm.op, src_dtypes, norm.dtype)
+      direct.setdefault(key, []).append(opcode)
+    else:
+      structural.append((opcode, norm))
+
+  # pick best opcode for each direct pattern
+  direct_best = {k: _select_best_opcode(v) for k, v in direct.items()}
+
+  # deduplicate structural patterns by shape, pick best
+  seen: dict[str, list] = {}
+  for opcode, norm in structural:
+    key = _pattern_key(norm)
+    seen.setdefault(key, []).append((opcode, norm))
+  structural_best = []
+  for key, group in seen.items():
+    best = _select_best_opcode([oc for oc, _ in group])
+    best_norm = next(n for oc, n in group if oc == best)
+    structural_best.append((best, best_norm, key))
+
+  return direct_best, structural_best
+
+# ═══════════════════════════════════════════════════════════════
+# Global tables (populated by build_isel_patterns, used by callbacks)
+# ═══════════════════════════════════════════════════════════════
+
+# direct: (op, src_dtypes, dst_dtype) -> Inst
+_DIRECT_TABLE: dict[tuple, object] = {}
+# structural: pattern_key_string -> Inst
+_STRUCTURAL_TABLE: dict[str, object] = {}
+
+# ops that LLVM handles natively (compares write VCC, not VGPRs)
+_SKIP_OPS = frozenset({Ops.CMPLT, Ops.CMPNE, Ops.CMPEQ})
+
+def _isel_direct(m):
+  """Callback for direct ALU: look up Inst by (op, src_dtypes, dtype)."""
+  if m.op in _SKIP_OPS or m.dtype == dtypes.bool: return None
+  src_dtypes = tuple(s.dtype for s in m.src)
+  inst = _DIRECT_TABLE.get((m.op, src_dtypes, m.dtype))
+  if inst is None: return None
+  return UOp(Ops.INS, m.dtype, m.src, arg=inst)
+
+def _isel_structural(m, **kwargs):
+  """Callback for structural patterns: compute runtime key, look up Inst."""
+  if m.op in _SKIP_OPS or m.dtype == dtypes.bool: return None
+  key = _runtime_key(m)
+  inst = _STRUCTURAL_TABLE.get(key)
+  if inst is None: return None
+  # collect source vars in order (leaves of the matched tree)
+  srcs = _collect_leaves(m)
+  return UOp(Ops.INS, m.dtype, tuple(srcs), arg=inst)
+
+def _collect_leaves(uop, _seen=None):
+  """Collect leaf UOps (non-ALU) from a matched tree in left-to-right order."""
+  if _seen is None: _seen = set()
+  _ALU_OPS = GroupOp.ALU | {Ops.CAST, Ops.BITCAST, Ops.WHERE}
+  uid = id(uop)
+  if uid in _seen: return []
+  _seen.add(uid)
+  if uop.op in (Ops.CONST, Ops.VCONST): return []
+  if uop.op not in _ALU_OPS: return [uop]
+  result = []
+  for s in uop.src:
+    result.extend(_collect_leaves(s, _seen))
+  return result
+
+# ═══════════════════════════════════════════════════════════════
+# Build the PatternMatcher
+# ═══════════════════════════════════════════════════════════════
+
+def build_isel_patterns(pcode_dict=PCODE, vgpr_only=False) -> PatternMatcher:
+  """Parse pcode and build a PatternMatcher for instruction selection."""
+  direct_best, structural_best = _parse_pcode_patterns(pcode_dict, vgpr_only=vgpr_only)
+
+  # populate direct table
+  _DIRECT_TABLE.clear()
+  for (op, src_dtypes, dtype), opcode in direct_best.items():
+    _DIRECT_TABLE[(op, src_dtypes, dtype)] = make_inst(opcode)
+
+  # populate structural table
+  _STRUCTURAL_TABLE.clear()
+  for opcode, norm, pkey in structural_best:
+    _STRUCTURAL_TABLE[pkey] = make_inst(opcode)
+
+  patterns: list[tuple] = []
+
+  # structural patterns first (more specific, should match before catch-all direct)
+  for opcode, norm, pkey in structural_best:
+    pat = uop_to_upat(norm).named('m')
+    patterns.append((pat, _isel_structural))
+
+  # direct ALU: catch-all patterns that look up by (op, src_dtypes, dtype)
+  patterns.append((UPat(GroupOp.ALU, name='m'), _isel_direct))
+  patterns.append((UPat(Ops.CAST, name='m'), _isel_direct))
+  patterns.append((UPat(Ops.BITCAST, name='m'), _isel_direct))
+
+  return PatternMatcher(patterns)
+
+@functools.cache
+def rdna3_isel() -> PatternMatcher:
+  """Build the default RDNA3 instruction selector (VOP-only for LLVM inline asm)."""
+  return build_isel_patterns(PCODE, vgpr_only=True)
+
+# ═══════════════════════════════════════════════════════════════
+# LLVM inline asm rendering for Ops.INS
+# ═══════════════════════════════════════════════════════════════
+
+import re
+from tinygrad.dtype import PtrDType
+
+def _ins_mnemonic(inst) -> str:
+  """Get the assembly mnemonic from an Inst opcode (strip _E32/_E64 suffix)."""
+  return re.sub(r'_e(32|64)$', '', inst.op.name.lower())
+
+def _ldt(dt):
+  """LLVM type string for a DType."""
+  if dt.vcount > 1: return f"<{dt.vcount} x {_ldt(dt.scalar())}>"
+  if isinstance(dt, PtrDType): return _ldt(dt.base) + "*"
+  return {dtypes.void: "void", dtypes.bool: "i1", dtypes.int8: "i8", dtypes.int16: "i16", dtypes.int32: "i32", dtypes.int64: "i64",
+          dtypes.uint8: "i8", dtypes.uint16: "i16", dtypes.uint32: "i32", dtypes.uint64: "i64",
+          dtypes.float16: "half", dtypes.bfloat16: "bfloat", dtypes.float32: "float", dtypes.float64: "double"}[dt]
+
+def render_ins_llvm(ctx, x):
+  """Render Ops.INS as LLVM inline assembly call."""
+  inst = x.arg
+  mnem = _ins_mnemonic(inst)
+  n_srcs = len(x.src)
+  # build operand string: $0 = dest, $1..$N = sources
+  ops = ", ".join(f"${i}" for i in range(n_srcs + 1))
+  asm_str = f"{mnem} {ops}"
+  # constraints: =v for output, v for each input (VGPR)
+  constraints = "=v," + ",".join("v" for _ in range(n_srcs))
+  # LLVM types and values
+  ret_type = _ldt(x.dtype)
+  args = ", ".join(f"{_ldt(s.dtype)} {ctx[s]}" for s in x.src)
+  return f"  {ctx[x]} = call {ret_type} asm \"{asm_str}\", \"{constraints}\"({args})"
+
+# ═══════════════════════════════════════════════════════════════
+# AMDISELRenderer: LLVM renderer with pcode-based instruction selection
+# ═══════════════════════════════════════════════════════════════
+
+from tinygrad.renderer.llvmir import AMDLLVMRenderer
+
+class AMDISELRenderer(AMDLLVMRenderer):
+  """AMD renderer that uses pcode-derived instruction selection for ALU ops."""
+  def __init__(self, arch: str):
+    super().__init__(arch)
+    # add ISel as extra_matcher: rewrites ALU UOps → Ops.INS
+    self.extra_matcher = self.extra_matcher + rdna3_isel()
+    # add Ops.INS rendering to string_rewrite
+    self.string_rewrite = PatternMatcher([(UPat(Ops.INS, name='x'), render_ins_llvm)]) + self.string_rewrite
+  def __reduce__(self): return self.__class__, (self.arch,)

tinygrad/renderer/amd/pcode.py

@@ -40,7 +40,10 @@ def _bitreverse(v: UOp, bits: int) -> UOp:
 
 def _extract_bits(val: UOp, hi: int, lo: int) -> UOp:
   dt = dtypes.uint64 if val.dtype in (dtypes.uint64, dtypes.int64) else dtypes.uint32
-  return ((val >> _const(dt, lo)) if lo > 0 else val) & _const(val.dtype, (1 << (hi - lo + 1)) - 1)
+  result = ((val >> _const(dt, lo)) if lo > 0 else val) & _const(val.dtype, (1 << (hi - lo + 1)) - 1)
+  # Downcast to uint32 when extracting <=32 bits from a 64-bit value, so .f32 bitcast works correctly
+  if dt == dtypes.uint64 and (hi - lo + 1) <= 32: result = result.cast(dtypes.uint32)
+  return result
 
 def _set_bit(old, pos, val):
   mask = _u32(1) << pos
@@ -554,7 +557,9 @@ class Parser:
       self.eat('LBRACKET')
       self.eat_val('laneId', 'IDENT')
       self.eat('RBRACKET')
-      result = (base >> _to_u32(self.vars['laneId'])) & _u32(1)
+      lane = self.vars['laneId']
+      shift = lane.cast(base.dtype) if base.dtype != dtypes.uint32 else _to_u32(lane)
+      result = (base >> shift) & _const(base.dtype, 1)
       if self.try_eat('DOT'):
         dt_name = self.eat('IDENT').val
         return result.cast(DTYPES.get(dt_name, dtypes.uint32))
@@ -806,6 +811,12 @@ def _subst_loop_var(line: str, loop_var: str, val: int) -> str:
 
 def _set_bits(old: UOp, val: UOp, width: int, offset: int) -> UOp:
   """Set bits [offset:offset+width) in old to val, masking and shifting appropriately."""
+  is64 = old.dtype in (dtypes.uint64, dtypes.int64) or offset + width > 32
+  if is64:
+    old = old.cast(dtypes.uint64) if old.dtype != dtypes.uint64 else old
+    mask = _u64(((1 << width) - 1) << offset)
+    v = (val.cast(dtypes.uint64) if val.dtype != dtypes.uint64 else val) & _u64((1 << width) - 1)
+    return (old & (mask ^ _u64(0xFFFFFFFFFFFFFFFF))) | (v << _u64(offset))
   mask = _u32(((1 << width) - 1) << offset)
   v = (val.cast(dtypes.uint32) if val.dtype != dtypes.uint32 else val) & _u32((1 << width) - 1)
   return (old & (mask ^ _u32(0xFFFFFFFF))) | (v << _u32(offset))

tinygrad/renderer/amd/renderer.py

@@ -0,0 +1,427 @@
+# Direct AMD GPU assembly renderer — emits Inst objects, produces GAS text via disasm()
+# No LLVM. Uses HIPCompiler (COMGR) to assemble text into ELF.
+
+import functools, math
+from tinygrad.uop.ops import UOp, UPat, Ops, PatternMatcher, GroupOp
+from tinygrad.dtype import dtypes, DType, PtrDType
+from tinygrad.renderer import Renderer
+from tinygrad.renderer.cstyle import AMDHIPRenderer
+from extra.assembly.amd.dsl import Inst, Reg, s, v, NULL, VCC_LO, EXEC_LO, M0
+from extra.assembly.amd.autogen.rdna3.ins import (s_load_b64, s_load_b128, s_mov_b32, s_waitcnt, s_endpgm, s_barrier,
+  s_branch, s_cbranch_scc0, s_cbranch_scc1, s_cmp_ge_i32, s_add_i32, s_and_b32, s_lshl_b32,
+  v_mov_b32_e32, v_add_f32_e32, v_add_nc_u32_e32, v_lshlrev_b32_e32, v_lshrrev_b32_e32,
+  v_and_b32_e32, v_mul_lo_u32, v_cmp_lt_i32_e32,
+  global_load_b32, global_load_b64, global_load_b128, global_store_b32, global_store_b64, global_store_b128,
+  ds_load_b32, ds_store_b32)
+from extra.assembly.amd.test.disasm import disasm
+from extra.assembly.amd.isel import rdna3_isel, make_inst
+
+# ═══════════════════════════════════════════════════════════════
+# Register allocator — simple bump allocator
+# ═══════════════════════════════════════════════════════════════
+
+class RegFile:
+  """Simple register allocator: bump-allocates VGPRs and SGPRs."""
+  def __init__(self):
+    self.next_vgpr = 1  # v0 = workitem_id_x (reserved by hardware)
+    self.next_sgpr = 0  # s[0:1] = kernarg_ptr (reserved by ABI)
+    self.max_vgpr = 1
+    self.max_sgpr = 0
+
+  def alloc_vgpr(self, count=1) -> Reg:
+    r = v[self.next_vgpr] if count == 1 else v[self.next_vgpr:self.next_vgpr + count - 1]
+    self.next_vgpr += count
+    self.max_vgpr = max(self.max_vgpr, self.next_vgpr)
+    return r
+
+  def alloc_sgpr(self, count=1) -> Reg:
+    # align to 2 for 64-bit, 4 for 128-bit
+    if count >= 4: self.next_sgpr = (self.next_sgpr + 3) & ~3
+    elif count >= 2: self.next_sgpr = (self.next_sgpr + 1) & ~1
+    r = s[self.next_sgpr] if count == 1 else s[self.next_sgpr:self.next_sgpr + count - 1]
+    self.next_sgpr += count
+    self.max_sgpr = max(self.max_sgpr, self.next_sgpr)
+    return r
+
+# ═══════════════════════════════════════════════════════════════
+# Instruction emitter (like amd_asm_matmul.Kernel)
+# ═══════════════════════════════════════════════════════════════
+
+class AsmKernel:
+  def __init__(self, arch='gfx1100'):
+    self.instructions: list[Inst] = []
+    self.labels: dict[str, int] = {}
+    self.pos = 0
+    self.arch = arch
+    self.regs = RegFile()
+    self.lds_size = 0
+
+  def emit(self, inst, target=None):
+    self.instructions.append(inst)
+    inst._target = target
+    inst._pos = self.pos
+    self.pos += inst.size()
+    return inst
+
+  def label(self, name):
+    self.labels[name] = self.pos
+
+  def waitcnt(self, lgkm=None, vm=None):
+    vmcnt = vm if vm is not None else 63
+    lgkmcnt = lgkm if lgkm is not None else 63
+    expcnt = 7
+    wc = (expcnt & 0x7) | ((lgkmcnt & 0x3f) << 4) | ((vmcnt & 0x3f) << 10)
+    self.emit(s_waitcnt(simm16=wc))
+
+  def resolve_branches(self):
+    for inst in self.instructions:
+      if hasattr(inst, '_target') and inst._target is not None:
+        offset_dwords = (self.labels[inst._target] - inst._pos - inst.size()) // 4
+        inst.simm16 = offset_dwords
+
+  def to_asm(self, name='kernel', kernarg_size=0, n_params=0) -> str:
+    self.resolve_branches()
+    body = ['\t' + disasm(inst) for inst in self.instructions]
+
+    hsa = [
+      ('group_segment_fixed_size', self.lds_size), ('private_segment_fixed_size', 0), ('kernarg_size', kernarg_size),
+      ('user_sgpr_count', 2), ('user_sgpr_kernarg_segment_ptr', 1),
+      ('wavefront_size32', 1), ('uses_dynamic_stack', 0), ('enable_private_segment', 0),
+      ('system_sgpr_workgroup_id_x', 1), ('system_sgpr_workgroup_id_y', 1), ('system_sgpr_workgroup_id_z', 0),
+      ('system_vgpr_workitem_id', 0), ('next_free_vgpr', self.regs.max_vgpr),
+      ('next_free_sgpr', max(self.regs.max_sgpr, 4)),  # minimum 4 SGPRs
+      ('float_round_mode_32', 0), ('float_round_mode_16_64', 0),
+      ('float_denorm_mode_32', 3), ('float_denorm_mode_16_64', 3),
+      ('dx10_clamp', 1), ('ieee_mode', 1), ('fp16_overflow', 0),
+      ('workgroup_processor_mode', 0), ('memory_ordered', 1), ('forward_progress', 0), ('shared_vgpr_count', 0)]
+
+    args_meta = '\n'.join(
+      f'      - .address_space: global\n        .offset: {i*8}\n        .size: 8\n        .value_kind: global_buffer'
+      for i in range(n_params))
+
+    return '\n'.join([
+      '\t.text', f'\t.amdgcn_target "amdgcn-amd-amdhsa--{self.arch}"',
+      f'\t.protected\t{name}', f'\t.globl\t{name}', '\t.p2align\t8', f'\t.type\t{name},@function', f'{name}:',
+      *body,
+      '\t.section\t.rodata,"a",@progbits', '\t.p2align\t6, 0x0', f'\t.amdhsa_kernel {name}',
+      *[f'\t\t.amdhsa_{k} {v}' for k, v in hsa],
+      f'\t.end_amdhsa_kernel', '\t.text', f'.Lfunc_end0:', f'\t.size\t{name}, .Lfunc_end0-{name}',
+      '\t.amdgpu_metadata', '---', 'amdhsa.kernels:', '  - .args:',
+      args_meta,
+      f'    .group_segment_fixed_size: {self.lds_size}', '    .kernarg_segment_align: 8',
+      f'    .kernarg_segment_size: {kernarg_size}', '    .max_flat_workgroup_size: 1024',
+      f'    .name: {name}', '    .private_segment_fixed_size: 0',
+      f'    .sgpr_count: {max(self.regs.max_sgpr, 4)}', f'    .symbol: {name}.kd',
+      f'    .vgpr_count: {self.regs.max_vgpr}', '    .wavefront_size: 32',
+      f'amdhsa.target: amdgcn-amd-amdhsa--{self.arch}',
+      'amdhsa.version:', '  - 1', '  - 2', '...', '\t.end_amdgpu_metadata'])
+
+# ═══════════════════════════════════════════════════════════════
+# UOp → Inst rendering
+# ═══════════════════════════════════════════════════════════════
+
+# dtype → register count for VGPRs
+def _dtype_regs(dt: DType) -> int:
+  if isinstance(dt, PtrDType): return 2  # 64-bit pointer
+  return max(1, dt.itemsize // 4) * (dt.vcount if hasattr(dt, 'vcount') and dt.vcount > 1 else 1)
+
+# dtype → global load instruction
+def _global_load(vdst, addr, saddr, offset=0, nregs=1):
+  if nregs == 1: return global_load_b32(vdst=vdst, addr=addr, saddr=saddr, offset=offset)
+  if nregs == 2: return global_load_b64(vdst=vdst, addr=addr, saddr=saddr, offset=offset)
+  if nregs == 4: return global_load_b128(vdst=vdst, addr=addr, saddr=saddr, offset=offset)
+  raise RuntimeError(f"unsupported global load size: {nregs} regs")
+
+def _global_store(addr, data, saddr, offset=0, nregs=1):
+  if nregs == 1: return global_store_b32(addr=addr, data=data, saddr=saddr, offset=offset)
+  if nregs == 2: return global_store_b64(addr=addr, data=data, saddr=saddr, offset=offset)
+  if nregs == 4: return global_store_b128(addr=addr, data=data, saddr=saddr, offset=offset)
+  raise RuntimeError(f"unsupported global store size: {nregs} regs")
+
+def render_kernel(uops: list[UOp], arch='gfx1100') -> str:
+  """Render linearized UOps into GAS assembly text."""
+  k = AsmKernel(arch)
+  # r maps UOp → register (Reg)
+  r: dict[UOp, Reg] = {}
+  # s_args: SGPR pairs for kernel argument pointers, loaded from kernarg segment
+  # kernarg_ptr is in s[0:1] (set by HSA ABI)
+  kernarg_base = k.regs.alloc_sgpr(2)  # s[0:1] = kernarg segment pointer
+  # system SGPRs for workgroup IDs come after user SGPRs
+  # with user_sgpr_count=2, workgroup_id_x is s[2], workgroup_id_y is s[3]
+  wg_id_x_sgpr = 2
+  wg_id_y_sgpr = 3
+
+  name = 'test'
+  params: list[tuple[int, Reg]] = []  # (param_idx, sgpr_pair)
+  specials: dict[str, Reg] = {}
+  loop_stack: list[tuple[str, str, Reg]] = []  # (label_start, label_end, range_reg)
+  n_params = 0
+
+  # first pass: count params
+  for u in uops:
+    if u.op is Ops.PARAM: n_params = max(n_params, u.arg + 1)
+    if u.op is Ops.SINK and u.arg is not None: name = u.arg.function_name
+
+  kernarg_size = n_params * 8  # each param is 8 bytes (pointer)
+
+  # load all kernel argument pointers
+  param_sgprs: dict[int, Reg] = {}
+  for i in range(n_params):
+    sp = k.regs.alloc_sgpr(2)
+    param_sgprs[i] = sp
+    k.emit(s_load_b64(sdata=sp, sbase=kernarg_base, offset=i * 8, soffset=NULL))
+  k.waitcnt(lgkm=0)
+
+  for u in uops:
+    if u.op is Ops.SINK:
+      continue
+
+    elif u.op is Ops.PARAM:
+      r[u] = param_sgprs[u.arg]
+
+    elif u.op is Ops.CONST:
+      if u.dtype == dtypes.float:
+        vr = k.regs.alloc_vgpr()
+        k.emit(v_mov_b32_e32(vr, u.arg))
+        r[u] = vr
+      elif u.dtype in (dtypes.int, dtypes.int32, dtypes.uint, dtypes.uint32):
+        vr = k.regs.alloc_vgpr()
+        k.emit(v_mov_b32_e32(vr, u.arg if isinstance(u.arg, int) and -16 <= u.arg <= 64 else u.arg))
+        r[u] = vr
+      elif u.dtype == dtypes.bool:
+        # booleans: 1=true, 0=false — stored in VGPR as int
+        vr = k.regs.alloc_vgpr()
+        k.emit(v_mov_b32_e32(vr, 1 if u.arg else 0))
+        r[u] = vr
+      else:
+        raise RuntimeError(f"unsupported CONST dtype {u.dtype}")
+
+    elif u.op is Ops.SPECIAL:
+      kind, idx = u.arg[0], int(u.arg[-1])
+      if kind == 'l':
+        # local thread ID — workitem_id_{x,y,z} already in v0 (only x for 1D)
+        if idx == 0:
+          r[u] = v[0]  # workitem_id_x is pre-loaded in v0 by hardware
+        else:
+          raise RuntimeError(f"unsupported local dim {idx}")
+      elif kind == 'g':
+        # workgroup ID — in system SGPRs (after user SGPRs)
+        sgpr_off = wg_id_x_sgpr + idx
+        vr = k.regs.alloc_vgpr()
+        k.emit(v_mov_b32_e32(vr, s[sgpr_off]))
+        r[u] = vr
+      else:
+        raise RuntimeError(f"unsupported SPECIAL kind {kind}")
+
+    elif u.op is Ops.INDEX:
+      # INDEX(ptr, idx) — compute byte address: base_ptr + idx * element_size
+      base = r[u.src[0]]
+      idx_reg = r[u.src[1]]
+      assert isinstance(u.dtype, PtrDType), f"INDEX must produce pointer, got {u.dtype}"
+      elem_size = u.dtype.base.itemsize
+      # compute byte offset: idx * elem_size
+      offset_vr = k.regs.alloc_vgpr()
+      if elem_size == 4:
+        k.emit(v_lshlrev_b32_e32(offset_vr, 2, idx_reg))
+      elif elem_size == 8:
+        k.emit(v_lshlrev_b32_e32(offset_vr, 3, idx_reg))
+      elif elem_size == 16:
+        k.emit(v_lshlrev_b32_e32(offset_vr, 4, idx_reg))
+      elif elem_size == 2:
+        k.emit(v_lshlrev_b32_e32(offset_vr, 1, idx_reg))
+      elif elem_size == 1:
+        k.emit(v_mov_b32_e32(offset_vr, idx_reg))
+      else:
+        k.emit(v_mul_lo_u32(offset_vr, elem_size, idx_reg))
+      # base is an SGPR pair (64-bit pointer), offset is VGPR — use scalar+vector addressing
+      r[u] = offset_vr  # store offset VGPR; base SGPR pair stored separately
+      # stash the base pointer for LOAD/STORE to use
+      u._base_sgpr = base
+
+    elif u.op is Ops.LOAD:
+      idx_uop = u.src[0]
+      assert idx_uop.op is Ops.INDEX or (idx_uop.op is Ops.CAST and idx_uop.src[0].op is Ops.INDEX)
+      real_idx = idx_uop.src[0] if idx_uop.op is Ops.CAST else idx_uop
+      base_sgpr = real_idx._base_sgpr
+      offset_vr = r[real_idx]
+      nregs = max(1, u.dtype.itemsize // 4) * (u.dtype.vcount if hasattr(u.dtype, 'vcount') and u.dtype.vcount > 1 else 1)
+      dst = k.regs.alloc_vgpr(nregs)
+      k.emit(_global_load(dst, offset_vr, base_sgpr, nregs=nregs))
+      k.waitcnt(vm=0)
+      r[u] = dst
+
+    elif u.op is Ops.STORE:
+      idx_uop = u.src[0]
+      assert idx_uop.op is Ops.INDEX or (idx_uop.op is Ops.CAST and idx_uop.src[0].op is Ops.INDEX)
+      real_idx = idx_uop.src[0] if idx_uop.op is Ops.CAST else idx_uop
+      base_sgpr = real_idx._base_sgpr
+      offset_vr = r[real_idx]
+      val_reg = r[u.src[1]]
+      nregs = max(1, u.src[1].dtype.itemsize // 4) * (u.src[1].dtype.vcount if hasattr(u.src[1].dtype, 'vcount') and u.src[1].dtype.vcount > 1 else 1)
+      k.emit(_global_store(offset_vr, val_reg, base_sgpr, nregs=nregs))
+      r[u] = offset_vr  # stores don't produce values, but map for dependencies
+
+    elif u.op is Ops.ADD:
+      a_reg, b_reg = r[u.src[0]], r[u.src[1]]
+      dst = k.regs.alloc_vgpr()
+      if u.dtype == dtypes.float:
+        k.emit(v_add_f32_e32(dst, a_reg, b_reg))
+      elif u.dtype in (dtypes.int, dtypes.int32, dtypes.uint, dtypes.uint32):
+        k.emit(v_add_nc_u32_e32(dst, a_reg, b_reg))
+      else:
+        raise RuntimeError(f"unsupported ADD dtype {u.dtype}")
+      r[u] = dst
+
+    elif u.op is Ops.MUL:
+      a_reg, b_reg = r[u.src[0]], r[u.src[1]]
+      dst = k.regs.alloc_vgpr()
+      if u.dtype == dtypes.float:
+        from extra.assembly.amd.autogen.rdna3.ins import v_mul_f32_e32
+        k.emit(v_mul_f32_e32(dst, a_reg, b_reg))
+      elif u.dtype in (dtypes.int, dtypes.int32, dtypes.uint, dtypes.uint32):
+        k.emit(v_mul_lo_u32(dst, a_reg, b_reg))
+      else:
+        raise RuntimeError(f"unsupported MUL dtype {u.dtype}")
+      r[u] = dst
+
+    elif u.op is Ops.SHL:
+      # SHL(val, shift) -> v_lshlrev_b32(shift, val) (reversed operands)
+      val_reg, shift_reg = r[u.src[0]], r[u.src[1]]
+      dst = k.regs.alloc_vgpr()
+      k.emit(v_lshlrev_b32_e32(dst, shift_reg, val_reg))
+      r[u] = dst
+
+    elif u.op is Ops.SHR:
+      val_reg, shift_reg = r[u.src[0]], r[u.src[1]]
+      dst = k.regs.alloc_vgpr()
+      k.emit(v_lshrrev_b32_e32(dst, shift_reg, val_reg))
+      r[u] = dst
+
+    elif u.op is Ops.AND:
+      a_reg, b_reg = r[u.src[0]], r[u.src[1]]
+      dst = k.regs.alloc_vgpr()
+      k.emit(v_and_b32_e32(dst, a_reg, b_reg))
+      r[u] = dst
+
+    elif u.op is Ops.CAST:
+      # for now: pointer casts are noops, numeric casts need work
+      if isinstance(u.dtype, PtrDType):
+        r[u] = r[u.src[0]]
+        if hasattr(u.src[0], '_base_sgpr'): u._base_sgpr = u.src[0]._base_sgpr
+      else:
+        raise RuntimeError(f"unsupported CAST {u.src[0].dtype} -> {u.dtype}")
+
+    elif u.op is Ops.VECTORIZE:
+      # VECTORIZE packs scalars into a vector — just allocate contiguous VGPRs
+      count = len(u.src)
+      dst = k.regs.alloc_vgpr(count)
+      for i, src_u in enumerate(u.src):
+        src_reg = r[src_u]
+        target = v[dst.offset - 256 + i] if count > 1 else dst
+        if src_reg.offset != target.offset:
+          k.emit(v_mov_b32_e32(target, src_reg))
+      r[u] = dst
+
+    elif u.op is Ops.GEP:
+      # GEP extracts element from vector — just offset into the VGPR range
+      base_reg = r[u.src[0]]
+      idx = u.arg[0]
+      r[u] = v[base_reg.offset - 256 + idx]
+
+    elif u.op in (Ops.NOOP, Ops.GROUP, Ops.AFTER):
+      if u.src: r[u] = r[u.src[0]]
+
+    elif u.op is Ops.RANGE:
+      # loop: counter starts at 0, increments by 1, bound is src[0]
+      label_start = f'loop_{id(u)}'
+      label_end = f'end_{id(u)}'
+      ctr = k.regs.alloc_vgpr()
+      k.emit(v_mov_b32_e32(ctr, 0))
+      k.label(label_start)
+      r[u] = ctr
+      loop_stack.append((label_start, label_end, ctr))
+
+    elif u.op is Ops.END:
+      label_start, label_end, ctr = loop_stack.pop()
+      # increment counter
+      k.emit(v_add_nc_u32_e32(ctr, 1, ctr))
+      # compare and branch: use SGPR compare since loop bound should be uniform
+      bound_uop = u.src[1]  # the RANGE uop's src[0] is the bound
+      # actually END.src = (range_uop, ...), range_uop.src[0] = bound
+      range_uop = u.src[0]
+      bound_reg = r[range_uop.src[0]]
+      k.emit(v_cmp_lt_i32_e32(ctr, bound_reg))
+      k.emit(s_cbranch_scc1(), target=label_start)
+      k.label(label_end)
+
+    elif u.op is Ops.BARRIER:
+      k.emit(s_barrier())
+
+    elif u.op is Ops.DEFINE_LOCAL:
+      # LDS allocation — just track size, address computed at use time
+      r[u] = v[0]  # placeholder, LDS addressing handled separately
+      k.lds_size = max(k.lds_size, u.dtype.size * u.dtype.base.itemsize if hasattr(u.dtype, 'size') else 0)
+
+    elif u.op is Ops.DEFINE_REG:
+      # register "spill" region — allocate VGPRs
+      size = u.dtype.size if hasattr(u.dtype, 'size') else 1
+      vr = k.regs.alloc_vgpr(size)
+      r[u] = vr
+
+    elif u.op is Ops.CMPLT:
+      # compare: write result to VCC, then v_cndmask to get bool in VGPR
+      a_reg, b_reg = r[u.src[0]], r[u.src[1]]
+      dst = k.regs.alloc_vgpr()
+      k.emit(v_cmp_lt_i32_e32(a_reg, b_reg))
+      from extra.assembly.amd.autogen.rdna3.ins import v_cndmask_b32_e32
+      k.emit(v_cndmask_b32_e32(dst, 0, 1))
+      r[u] = dst
+
+    elif u.op is Ops.CMPNE:
+      from extra.assembly.amd.autogen.rdna3.ins import v_cmp_ne_u32_e32, v_cndmask_b32_e32
+      a_reg, b_reg = r[u.src[0]], r[u.src[1]]
+      dst = k.regs.alloc_vgpr()
+      k.emit(v_cmp_ne_u32_e32(a_reg, b_reg))
+      k.emit(v_cndmask_b32_e32(dst, 0, 1))
+      r[u] = dst
+
+    elif u.op is Ops.WHERE:
+      from extra.assembly.amd.autogen.rdna3.ins import v_cndmask_b32_e32
+      cond_reg, true_reg, false_reg = r[u.src[0]], r[u.src[1]], r[u.src[2]]
+      dst = k.regs.alloc_vgpr()
+      # set VCC from condition (nonzero = true)
+      k.emit(v_cmp_lt_i32_e32(0, cond_reg))  # VCC = cond_reg != 0
+      k.emit(v_cndmask_b32_e32(dst, false_reg, true_reg))
+      r[u] = dst
+
+    else:
+      raise RuntimeError(f"unsupported UOp: {u.op} dtype={u.dtype}")
+
+  # epilogue
+  k.waitcnt(vm=0, lgkm=0)
+  k.emit(s_endpgm())
+
+  return k.to_asm(name=name, kernarg_size=kernarg_size, n_params=n_params)
+
+# ═══════════════════════════════════════════════════════════════
+# AMDAssemblyRenderer: Renderer subclass for tinygrad integration
+# ═══════════════════════════════════════════════════════════════
+
+class AMDAssemblyRenderer(Renderer):
+  device = "AMD"
+  suffix = "s"  # GAS assembly
+  supports_float4 = True
+  has_local = True
+  has_shared = True
+  global_max = AMDHIPRenderer.global_max
+  shared_max = AMDHIPRenderer.shared_max
+
+  def __init__(self, arch: str):
+    from tinygrad.runtime.support.compiler_amd import HIPCompiler
+    self.arch = arch
+    self.compiler = HIPCompiler(arch)
+
+  def render(self, uops: list[UOp]) -> str:
+    return render_kernel(uops, arch=self.arch)
+
+  def __reduce__(self): return self.__class__, (self.arch,)

tinygrad/renderer/amd/sqtt.py

@@ -670,7 +670,8 @@ if __name__ == "__main__":
     sys.exit(1)
   with open(sys.argv[1], "rb") as f:
     data = pickle.load(f)
+  prg_names = {e.tag: e.name for e in data if type(e).__name__ == "ProfileProgramEvent" and e.tag is not None}
   sqtt_events = [e for e in data if type(e).__name__ == "ProfileSQTTEvent"]
   for i, event in enumerate(sqtt_events):
-    print(f"\n=== event {i} ===")
+    print(f"\n=== event {i} {prg_names.get(event.kern, '')} ===")
     print_packets(decode(event.blob))

tinygrad/runtime/ops_amd.py

@@ -8,9 +8,11 @@ from tinygrad.runtime.support.hcq import MMIOInterface, BumpAllocator, hcq_filte
 from tinygrad.uop.ops import sint
 from tinygrad.device import Compiled, DMAFdRef, BufferSpec, CompilerSet
 from tinygrad.helpers import getenv, round_up, data64_le, DEBUG, PROFILE, ProfileEvent, lo32, hi32, colored, prod, ContextVar
-from tinygrad.helpers import VIZ, AMD_CC, AMD_LLVM, AMD_HIPCC, ceildiv, unwrap
+from tinygrad.helpers import VIZ, AMD_CC, AMD_LLVM, AMD_HIPCC, AMD_ISEL, AMD_ASM, ceildiv, unwrap
 from tinygrad.renderer.cstyle import AMDHIPRenderer, AMDHIPCCRenderer
 from tinygrad.renderer.llvmir import AMDLLVMRenderer
+from extra.assembly.amd.isel import AMDISELRenderer
+from extra.assembly.amd.renderer import AMDAssemblyRenderer
 from tinygrad.runtime.autogen import kfd, hsa, pci, sqtt, amdgpu_kd, amdgpu_drm
 from tinygrad.runtime.autogen.am import am
 from tinygrad.runtime.support.elf import elf_loader
@@ -970,7 +972,9 @@ class AMDDevice(HCQCompiled):
 
     compilers = CompilerSet([(functools.partial(AMDHIPRenderer, self.arch), None),
                              (functools.partial(AMDLLVMRenderer, self.arch), AMD_LLVM),
-                             (functools.partial(AMDHIPCCRenderer, self.arch), AMD_HIPCC)], ctrl_var=AMD_CC)
+                             (functools.partial(AMDHIPCCRenderer, self.arch), AMD_HIPCC),
+                             (functools.partial(AMDISELRenderer, self.arch), AMD_ISEL),
+                             (functools.partial(AMDAssemblyRenderer, self.arch), AMD_ASM)], ctrl_var=AMD_CC)
 
     super().__init__(device, AMDAllocator(self), compilers, functools.partial(AMDProgram, self), AMDSignal,
                      functools.partial(AMDComputeAQLQueue if self.is_aql else AMDComputeQueue, self),

tinygrad/runtime/support/am/amdev.py

@@ -193,7 +193,7 @@ class AMDev(PCIDevImplBase):
     if DEBUG >= 2: print(f"am {self.devfmt}: boot done")
 
   def init_sw(self, smi_dev=False):
-    self.smi_dev, self.is_err_state = smi_dev, False
+    self.smi_dev, self.is_err_state, self.has_aql_queue = smi_dev, False, False
 
     # Memory manager & firmware
     self.mm = AMMemoryManager(self, self.vram_size - self.reserved_vram_size, boot_size=(32 << 20), pt_t=AMPageTableEntry, va_shifts=[12, 21, 30, 39],
@@ -226,7 +226,7 @@ class AMDev(PCIDevImplBase):
     self.reg("regSCRATCH_REG6").write(self.is_err_state) # set finalized state.
 
   def recover(self) -> bool:
-    if self.is_hive() or not self.is_err_state: return False # TODO: support mi300
+    if (self.has_aql_queue and self.is_hive()) or not self.is_err_state: return False # TODO: support aql queue recovery on hive
     if DEBUG >= 2: print(f"am {self.devfmt}: Start recovery")
     self.ih.interrupt_handler()
     self.gfx.reset_mec()

tinygrad/runtime/support/am/ip.py

@@ -291,6 +291,7 @@ class AM_GFX(AM_IP):
     self._enable_mec()
 
   def setup_ring(self, ring_addr:int, ring_size:int, rptr_addr:int, wptr_addr:int, eop_addr:int, eop_size:int, idx:int, aql:bool) -> tuple[int, int]:
+    self.adev.has_aql_queue |= aql
     pipe, queue, doorbell = idx // 4, idx % 4, am.AMDGPU_NAVI10_DOORBELL_MEC_RING0
     self._grbm_select(me=1, pipe=pipe, queue=queue, inst=0)
     restore_queue = aql and self.xccs > 1 and self.adev.partial_boot and (self.adev.regCP_HQD_ACTIVE.read(inst=0) & 1)

tinygrad/uop/__init__.py

@@ -76,7 +76,7 @@ class Ops(FastEnum):
   # CUSTOM/CUSTOMI are used to output strings into codegen. the I makes the string inline
   CUSTOM = auto(); CUSTOMI = auto()
 
-  # INS is a machine instruction
+  # machine instruction: arg=Inst object, tag=register assignment
   INS = auto()
 
   # ** 6 -- ops that don't exist in programs **

tinygrad/uop/ops.py

@@ -289,7 +289,7 @@ class UOp(OpMixin, metaclass=UOpMetaClass):
     if self.op is Ops.ASSIGN: return self.src[1]._shape
 
     # elementwise ops keep the shape the same. all inputs with shape must match
-    if self.op in GroupOp.ALU.union({Ops.CAST, Ops.COPY, Ops.NOOP, Ops.GROUP, Ops.SINK, Ops.ALLREDUCE, Ops.STORE}):
+    if self.op in GroupOp.ALU.union({Ops.CAST, Ops.COPY, Ops.NOOP, Ops.GROUP, Ops.SINK, Ops.ALLREDUCE, Ops.STORE, Ops.INS}):
       input_shapes = [x._shape for x in self.src if x._shape is not None]
       if len(input_shapes) == 0: return None
       if not all_same(input_shapes): raise RuntimeError(f"shape mismatch at {self.op}: {input_shapes}")

tinygrad/uop/spec.py

@@ -177,7 +177,7 @@ shared_codegen_spec = PatternMatcher([
   # CUSTOM (inline and non inline)
   (UPat((Ops.CUSTOMI, Ops.CUSTOM)), lambda: True),
 
-  # assembly instruction
+  # machine instruction (ISel output)
   (UPat(Ops.INS), lambda: True),
 
   # INDEX (2-arg and 3-arg with bool gate)

tinygrad/viz/js/index.js

@@ -184,7 +184,7 @@ const WAVE_COLORS = {VALU:"#ffffc0", SALU:"#cef263", LOAD:"#ffc0c0", STORE:"#4fa
 const waveColor = (op) => {
   const cat = op.includes("VALU") || op === "VINTERP" ? "VALU" : op.includes("SALU") ? "SALU" : op.includes("VMEM") ? "VMEM"
             : op.includes("LOAD") || op === "SMEM" ? "LOAD" : op.includes("STORE") ? "STORE" : op;
-  ret = WAVE_COLORS[cat] ?? "#ffffff";
+  let ret = WAVE_COLORS[cat] ?? "#ffffff";
   if (op.includes("OTHER_") || op.includes("_ALT")) { ret = darkenHex(ret, 75) }
   if (op.includes("LDS_")) { ret = darkenHex(ret, 25) }
   return ret