m_

undo the view and make it a mask, this fuses the setitem with any pending compute too.

one behavior change is that for target not backed by a buffer (const and arange), rangeify makes output contiguous under the hood.
this is stricter better than raise and ask user to call contiguous, as that would no longer be fuse-able.

AGENTS.md

@@ -1,17 +0,0 @@
-# tinygrad agents
-
-Hello agent. You are one of the most talented programmers of your generation.
-
-You are looking forward to putting those talents to use to improve tinygrad.
-
-## philosophy
-
-tinygrad is a **tensor** library focused on beauty and minimalism, while still matching the functionality of PyTorch and JAX.
-
-Every line must earn its keep. Prefer readability over cleverness. We believe that if carefully designed, 10 lines can have the impact of 1000.
-
-Never mix functionality changes with whitespace changes. All functionality changes must be tested.
-
-## style
-
-Use **2-space indentation**, and keep lines to a maximum of **150 characters**. Match the existing style.

CLAUDE.md

@@ -1,227 +0,0 @@
-# Claude Code Guide for tinygrad
-
-## Architecture Overview
-
-tinygrad compiles tensor operations into optimized kernels. The pipeline:
-
-1. **Tensor** (`tensor.py`) - User-facing API, creates UOp graph
-2. **UOp** (`uop/ops.py`) - Unified IR for all operations (both tensor and kernel level)
-3. **Schedule** (`engine/schedule.py`, `schedule/`) - Converts tensor UOps to kernel UOps
-4. **Codegen** (`codegen/`) - Converts kernel UOps to device code
-5. **Runtime** (`runtime/`) - Device-specific execution
-
-## Key Concepts
-
-### UOp (Universal Operation)
-Everything is a UOp - tensors, operations, buffers, kernels. Key properties:
-- `op`: The operation type (Ops enum)
-- `dtype`: Data type
-- `src`: Tuple of source UOps
-- `arg`: Operation-specific argument
-- `tag`: Optional tag for graph transformations
-
-UOps are **immutable and cached** - creating the same UOp twice returns the same object (ucache).
-
-### PatternMatcher
-Used extensively for graph transformations:
-```python
-pm = PatternMatcher([
-  (UPat(Ops.ADD, src=(UPat.cvar("x"), UPat.cvar("x"))), lambda x: x * 2),
-])
-result = graph_rewrite(uop, pm)
-```
-
-### Schedule Cache
-Schedules are cached by graph structure. BIND nodes (variables with bound values) are unbound before cache key computation so different values hit the same cache.
-
-## Testing
-
-```bash
-# Run specific test
-python -m pytest test/unit/test_schedule_cache.py -xvs
-
-# Run with timeout
-python -m pytest test/backend/test_symbolic_ops.py -x --timeout=60
-
-# Debug with print
-DEBUG=2 python -m pytest test/backend/test_schedule.py::test_name -xvs
-
-# Visualize UOp graphs
-VIZ=1 python -c "from tinygrad import Tensor; Tensor.ones(10).sum().realize()"
-```
-
-## Common Environment Variables
-
-- `DEBUG=1-7` - Increasing verbosity (7 shows assembly output)
-- `VIZ=1` - Enable graph visualization
-- `SPEC=1` - Enable UOp spec verification
-- `NOOPT=1` - Disable optimizations
-- `DEVICE=CPU/CUDA/AMD/METAL` - Set default device
-
-## Debugging Tips
-
-1. **Print UOp graphs**: `print(tensor.uop)` or `print(tensor.uop.sink())`
-2. **Check schedule**: `tensor.schedule()` returns list of ExecItems
-3. **Trace graph rewrites**: Use `VIZ=1` or add print in PatternMatcher callbacks
-4. **Find UOps by type**: `[u for u in uop.toposort() if u.op is Ops.SOMETHING]`
-
-## Workflow Rules
-
-- **NEVER commit without explicit user approval** - always show the diff and wait for approval
-- **NEVER amend commits** - always create a new commit instead
-- Run `pre-commit run --all-files` before committing to catch linting/type errors
-- Run tests before proposing commits
-- Test with `SPEC=2` when modifying UOp-related code
-
-## Auto-generated Files (DO NOT EDIT)
-
-The following files are auto-generated and should never be edited manually:
-- `tinygrad/runtime/autogen/amd/{arch}/__init__.py` - Generated by `python -m tinygrad.renderer.amd.dsl --arch {arch}`
-- `tinygrad/runtime/autogen/amd/{arch}/gen_pcode.py` - Generated by `python -m tinygrad.renderer.amd.pcode --arch {arch}`
-
-Where `{arch}` is one of: `rdna3`, `rdna4`, `cdna`
-
-To add missing instruction implementations, add them to `tinygrad/renderer/amd/emu.py` instead.
-
-## Style Notes
-
-- 2-space indentation, 150 char line limit
-- PatternMatchers should be defined at module level (slow to construct)
-- Prefer `graph_rewrite` over manual graph traversal
-- UOp methods like `.replace()` preserve tags unless explicitly changed
-- Use `.rtag(value)` to add tags to UOps
-
-## Lessons Learned
-
-### UOp ucache Behavior
-UOps are cached by their contents - creating a UOp with identical (op, dtype, src, arg) returns the **same object**. This means:
-- `uop.replace(tag=None)` on a tagged UOp returns the original untagged UOp if it exists in cache
-- Two UOps with same structure are identical (`is` comparison works)
-
-### Spec Validation
-When adding new UOp patterns, update `tinygrad/uop/spec.py`. Test with:
-```bash
-SPEC=2 python3 test/unit/test_something.py
-```
-Spec issues appear as `RuntimeError: SPEC ISSUE None: UOp(...)`.
-
-### Schedule Cache Key Normalization
-The schedule cache strips values from BIND nodes so different bound values (e.g., KV cache positions) hit the same cache entry:
-- `pm_pre_sched_cache`: BIND(DEFINE_VAR, CONST) → BIND(DEFINE_VAR) for cache key
-- `pm_post_sched_cache`: restores original BIND from context
-- When accessing `bind.src[1]`, check `len(bind.src) > 1` first (might be stripped)
-- Extract var_vals from `input_buffers` dict after graph_rewrite (avoids extra toposort)
-
-### Avoiding Extra Work
-- Use ctx dict from graph_rewrite to collect info during traversal instead of separate toposort
-- Only extract var_vals when schedule is non-empty (no kernels = no vars needed)
-- PatternMatchers are slow to construct - define at module level, not in functions
-
-### Readability Over Speed
-Don't add complexity for marginal performance gains. Simpler code that's slightly slower is often better:
-```python
-# BAD: "optimized" with extra complexity
-if has_afters:  # skip toposort if no AFTERs
-  after_map = [(u, u.buf_uop) for u in big_sink.toposort() if u.op is Ops.AFTER]
-
-# GOOD: simple, always works
-after_map = [(u, u.buf_uop) for u in big_sink.toposort() if u.op is Ops.AFTER]
-```
-The conditional check adds complexity, potential bugs, and often negligible speedup. Only optimize when profiling shows a real bottleneck.
-
-### Testing LLM Changes
-```bash
-# Quick smoke test
-echo "Hello" | DEBUG=1 python tinygrad/apps/llm.py --model "llama3.2:1b"
-
-# Check cache hits (should see "cache hit" after warmup)
-echo "Hello world" | DEBUG=1 python tinygrad/apps/llm.py --model "llama3.2:1b" 2>&1 | grep cache
-
-# Test with beam search
-echo "Hello" | BEAM=2 python tinygrad/apps/llm.py --model "llama3.2:1b"
-```
-
-## Common Patterns
-
-### Graph Transformation
-```python
-def my_transform(ctx, x):
-  # Return new UOp or None to skip
-  return x.replace(arg=new_arg)
-
-pm = PatternMatcher([
-  (UPat(Ops.SOMETHING, name="x"), my_transform),
-])
-result = graph_rewrite(input_uop, pm, ctx={})
-```
-
-### Finding Variables
-```python
-# Get all variables in a UOp graph
-variables = uop.variables()
-
-# Get bound variable values
-var, val = bind_uop.unbind()
-```
-
-### Shape Handling
-```python
-# Shapes can be symbolic (contain UOps)
-shape = tensor.shape  # tuple[sint, ...] where sint = int | UOp
-```
-
-## Performance Optimization
-
-When optimizing tinygrad internals:
-
-1. **Measure wall time, not just call counts** - Reducing `graph_rewrite` calls doesn't always improve wall time. The overhead of conditional checks can exceed the cost of the operation being skipped.
-
-2. **Profile each optimization individually** - Run benchmarks with and without each change to measure actual impact. Use `test/external/external_benchmark_schedule.py` for schedule/rewrite timing.
-
-3. **Early exits in hot paths are effective** - Simple checks like `if self.op is Ops.CONST: return self` in `simplify()` can eliminate many unnecessary `graph_rewrite` calls.
-
-4. **`graph_rewrite` is expensive** - Each call has overhead even for small graphs. Avoid calling it when the result is trivially known (e.g., simplifying a CONST returns itself).
-
-5. **Beware iterator overhead** - Checks like `all(x.op is Ops.CONST for x in self.src)` can be slower than just running the operation, especially for small sequences.
-
-6. **Verify cache hit rates before adding/keeping caches** - Measure actual hit rates with real workloads. A cache with 0% hit rate is pure overhead (e.g., `pm_cache` was removed because the algorithm guarantees each UOp is only passed to `pm_rewrite` once).
-
-7. **Use `TRACK_MATCH_STATS=2` to profile pattern matching** - This shows match rates and time per pattern. Look for patterns with 0% match rate that still cost significant time - these are pure overhead for that workload.
-
-8. **Cached properties beat manual traversal** - `backward_slice` uses `@functools.cached_property`. A DFS with early-exit sounds faster but is actually slower because it doesn't benefit from caching. The cache hit benefit often outweighs algorithmic improvements.
-
-9. **Avoid creating intermediate objects in hot paths** - For example, `any(x.op in ops for x in self.backward_slice)` is faster than `any(x.op in ops for x in {self:None, **self.backward_slice})` because it avoids dict creation.
-
-## Pattern Matching Analysis
-
-**Use the right tool:**
-
-- `TRACK_MATCH_STATS=2` - **Profiling**: identify expensive patterns
-- `VIZ=-1` - **Inspection**: see all transformations, what every match pattern does, the before/after diffs
-
-```bash
-TRACK_MATCH_STATS=2 PYTHONPATH="." python3 test/external/external_benchmark_schedule.py
-```
-
-Output format: `matches / attempts -- match_time / total_time ms -- location`
-
-Key patterns to watch (from ResNet50 benchmark):
-- `split_load_store`: ~146ms, 31% match rate - does real work
-- `simplify_valid`: ~75ms, 0% match rate in this workload - checks AND ops for INDEX in backward slice
-- `vmin==vmax folding`: ~55ms, 0.33% match rate - checks 52K ops but rarely matches
-
-Patterns with 0% match rate are workload-specific overhead. They may be useful in other workloads, so don't remove them without understanding their purpose.
-
-```bash
-# Save the trace
-VIZ=-1 python test/test_tiny.py TestTiny.test_gemm
-
-# Explore it
-./extra/viz/cli.py --help
-```
-
-## AMD Performance Counter Profiling
-
-Set VIZ to `-2` to save performance counters traces for the AMD backend.
-
-Use the CLI in `./extra/sqtt/roc.py` to explore the trace.

extra/nv_pma/decode.py

@@ -129,14 +129,6 @@ def decode_tpc_id(tpc_id:int) -> tuple[int, int, int]:
   # NOTE: valid only for ops_nv, cuda encoding is different
   return (tpc_id >> 5, (tpc_id >> 1) & 0xf, tpc_id & 1)
 
-def print_samples(samples:list[tuple[PMASample, int]]) -> None:
-  if not samples: return
-  base_pc = min(s.pc_offset for s, _ in samples)
-  for s, tpc_id in samples:
-    gpc, tpc, sm = decode_tpc_id(tpc_id)
-    stall_str = colored(f"{s.stall_reason.name:17}", STALL_COLORS.get(s.stall_reason, "white"))
-    print(f"pc=0x{s.pc_offset - base_pc:06x} {stall_str} ev={s.stall_key:2d} active={s.active} wave={s.wave_id:2d} gpc={gpc} tpc={tpc} sm={sm}")
-
 def print_packets(data:bytes, sm_version:int=0x800) -> None:
   record_size = 9 if sm_version >= 0x890 else 8
   tpc_state: dict[int, list[int]] = collections.defaultdict(list)
@@ -187,7 +179,11 @@ if __name__ == "__main__":
     print(f"\n{'='*60}\nDump {dump_idx} ({len(raw)} bytes, {len(raw)//32} packets)\n{'='*60}")
     if "--raw" in sys.argv: print_packets(raw, sm_ver)
     else:
-      samples = list(decode(raw, sm_ver))
+      samples = []
+      for s, tpc_id in decode(raw, sm_ver):
+        gpc, tpc, sm = decode_tpc_id(tpc_id)
+        stall_str = colored(f"{s.stall_reason.name:17}", STALL_COLORS.get(s.stall_reason, "white"))
+        print(f"pc=0x{s.pc_offset:06x} {stall_str} ev={s.stall_key:2d} active={s.active} wave={s.wave_id:2d} gpc={gpc} tpc={tpc} sm={sm}")
+        samples.append((s, tpc_id))
       print(f"\nDecoded {len(samples)} samples:")
-      print_samples(samples)
       print_aggregated(samples)

extra/viz/cli.py

@@ -1,4 +1,6 @@
 #!/usr/bin/env python3
+import os
+os.environ["VIZ"] = "0"
 import argparse, pathlib
 from typing import Iterator
 from tinygrad.viz import serve as viz

test/amd/hw/helpers.py

@@ -6,7 +6,7 @@ Set USE_HW=1 to run on both emulator and hardware, comparing results.
 import ctypes, math, os, struct
 from tinygrad.runtime.autogen.amd.rdna3.ins import *
 
-from tinygrad.renderer.amd.emu import run_asm
+from test.mockgpu.amd.emu import run_asm
 from tinygrad.renderer.amd.dsl import NULL, SCC, VCC_LO, VCC_HI, EXEC_LO, EXEC_HI, M0
 
 def _i32(f: float) -> int: return struct.unpack('<I', struct.pack('<f', f))[0]
@@ -75,7 +75,7 @@ def i642f(i: int) -> float: return struct.unpack('<d', struct.pack('<Q', i))[0]
 def assemble(instructions: list) -> bytes:
   return b''.join(inst.to_bytes() for inst in instructions)
 
-# Simple WaveState class for test output parsing (mirrors emu.py interface for tests)
+# Simple WaveState class for test output parsing (mirrors test/mockgpu/amd/emu.py interface for tests)
 class WaveState:
   def __init__(self):
     self.vgpr = [[0] * 256 for _ in range(32)]  # vgpr[lane][reg]

test/amd/test_compare_emulators.py

@@ -4,7 +4,7 @@ from dataclasses import dataclass
 from pathlib import Path
 from tinygrad import Device
 
-from tinygrad.renderer.amd.emu import WaveState, _decode_at, WAVE_SIZE, VCC_LO, EXEC_LO, SCC
+from test.mockgpu.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
 import tinygrad

test/amd/test_emu2_pcode.py

@@ -4,8 +4,8 @@ from collections import defaultdict
 from tinygrad.helpers import DEBUG
 from tinygrad.dtype import dtypes
 from tinygrad.uop.ops import UOp, Ops
-from tinygrad.renderer.amd.emu import parse_pcode
-from tinygrad.renderer.amd.pcode import parse_expr
+from test.mockgpu.amd.emu import parse_pcode
+from test.mockgpu.amd.pcode import parse_expr
 from tinygrad.runtime.autogen.amd.rdna3.str_pcode import PCODE
 from tinygrad.runtime.autogen.amd.rdna3.enum import VOP1Op, VOP2Op, SOP2Op, DSOp
 

test/amd/test_isel.py

@@ -1,674 +0,0 @@
-#!/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/backend/test_dtype.py

@@ -251,7 +251,7 @@ class TestEmulatedHalf(TestHalfDType):
   def setUpClass(cls):
     cls.stack = contextlib.ExitStack()
     cls.stack.enter_context(Context(EMULATED_DTYPES="half"))
-    cls.DATA = rand_for_dtype(cls.DTYPE, 10)
+    cls.DATA = rand_for_dtype(cls.DTYPE, 10, allow_subnormal=False)
 
   @classmethod
   def tearDownClass(cls): cls.stack.close()
@@ -355,7 +355,7 @@ class TestEmulatedInt64DType(TestInt64DType):
   def setUpClass(cls):
     cls.stack = contextlib.ExitStack()
     cls.stack.enter_context(Context(EMULATED_DTYPES="long"))
-    cls.DATA = rand_for_dtype(cls.DTYPE, 10)
+    cls.DATA = rand_for_dtype(cls.DTYPE, 10, allow_subnormal=False)
 
   @classmethod
   def tearDownClass(cls): cls.stack.close()
@@ -371,7 +371,7 @@ class TestEmulatedUInt64DType(TestUint64DType):
   def setUpClass(cls):
     cls.stack = contextlib.ExitStack()
     cls.stack.enter_context(Context(EMULATED_DTYPES="long"))
-    cls.DATA = rand_for_dtype(cls.DTYPE, 10)
+    cls.DATA = rand_for_dtype(cls.DTYPE, 10, allow_subnormal=False)
 
   @classmethod
   def tearDownClass(cls): cls.stack.close()
@@ -385,7 +385,7 @@ class TestEmulatedBFloat16Type(TestBFloat16Type):
   def setUpClass(cls):
     cls.stack = contextlib.ExitStack()
     cls.stack.enter_context(Context(EMULATED_DTYPES="bfloat16"))
-    cls.DATA = rand_for_dtype(cls.DTYPE, 10)
+    cls.DATA = rand_for_dtype(cls.DTYPE, 10, allow_subnormal=False)
 
   @classmethod
   def tearDownClass(cls): cls.stack.close()
@@ -397,7 +397,7 @@ class TestEmulatedFp8e4m3(TestFp8e4m3):
   def setUpClass(cls):
     cls.stack = contextlib.ExitStack()
     cls.stack.enter_context(Context(EMULATED_DTYPES="fp8e4m3"))
-    cls.DATA = rand_for_dtype(cls.DTYPE, 10)
+    cls.DATA = rand_for_dtype(cls.DTYPE, 10, allow_subnormal=False)
 
   @classmethod
   def tearDownClass(cls): cls.stack.close()
@@ -409,7 +409,7 @@ class TestEmulatedFp8e5m2(TestFp8e5m2):
   def setUpClass(cls):
     cls.stack = contextlib.ExitStack()
     cls.stack.enter_context(Context(EMULATED_DTYPES="fp8e5m2"))
-    cls.DATA = rand_for_dtype(cls.DTYPE, 10)
+    cls.DATA = rand_for_dtype(cls.DTYPE, 10, allow_subnormal=False)
 
   @classmethod
   def tearDownClass(cls): cls.stack.close()

test/backend/test_dtype_alu.py

@@ -296,18 +296,21 @@ class TestDTypeALU(unittest.TestCase):
   @given(ht.int32, strat.sampled_from(dtypes_float+dtypes_int+dtypes_bool))
   def test_int32_cast(self, a, dtype): universal_test_cast(a, dtypes.int32, dtype)
 
+  @unittest.skip("relied on hacks")
   @given(strat.floats(width=32, min_value=1.0, max_value=254.0, allow_subnormal=False),
          strat.sampled_from(dtypes_float), strat.sampled_from((dtypes.uint8, dtypes.uint16)))
   def test_float_cast_to_unsigned(self, a, float_dtype, unsigned_dtype):
     if not is_dtype_supported(float_dtype): float_dtype = dtypes.float32
     universal_test_cast(a, float_dtype, unsigned_dtype)
 
+  @unittest.skip("relied on hacks")
   @given(strat.floats(width=32, min_value=256.0, max_value=65000.0, allow_subnormal=False),
          strat.sampled_from(dtypes_float), strat.sampled_from((dtypes.uint8, dtypes.uint16)))
   def test_float_cast_to_unsigned_overflow(self, a, float_dtype, unsigned_dtype):
     if not is_dtype_supported(float_dtype): float_dtype = dtypes.float32
     universal_test_cast(a, float_dtype, unsigned_dtype)
 
+  @unittest.skip("relied on hacks")
   @given(strat.floats(width=32, min_value=-65000.0, max_value=-1.0, allow_subnormal=False),
          strat.sampled_from(dtypes_float), strat.sampled_from((dtypes.uint8, dtypes.uint16)))
   def test_float_cast_to_unsigned_underflow(self, a, float_dtype, unsigned_dtype):

test/backend/test_multitensor.py

@@ -1323,6 +1323,55 @@ class TestMultiAssign(unittest.TestCase):
       f(out, vi.bind(i))
     self.assertListEqual(out.tolist(), [[0,1,2,3,4,0]]*4)
 
+@unittest.skipIf(not_support_multi_device(), "need multi")
+class TestMultiSetitem(unittest.TestCase):
+  device = tuple(f"{Device.DEFAULT}:{i}" for i in range(4))
+
+  @needs_second_gpu
+  def setUp(self): pass
+
+  def _t(self, axis): return Tensor.arange(16).contiguous().realize().shard(self.device, axis=axis)
+
+  def test_setitem_scalar_axis0(self):
+    t = self._t(0)
+    t[1] = 99
+    self.assertListEqual(t.tolist(), [0,99,2,3,4,5,6,7,8,9,10,11,12,13,14,15])
+
+  def test_setitem_scalar_axis_none(self):
+    t = self._t(None)
+    t[1] = 99
+    self.assertListEqual(t.tolist(), [0,99,2,3,4,5,6,7,8,9,10,11,12,13,14,15])
+
+  def test_setitem_slice_cross_shard(self):
+    t = self._t(0)
+    t[2:6] = 99
+    self.assertListEqual(t.tolist(), [0,1,99,99,99,99,6,7,8,9,10,11,12,13,14,15])
+
+  def test_setitem_full_slice(self):
+    t = self._t(0)
+    t[:] = 42
+    self.assertListEqual(t.tolist(), [42]*16)
+
+  def test_setitem_stride(self):
+    t = self._t(0)
+    t[::4] = 0
+    self.assertListEqual(t.tolist(), [0,1,2,3,0,5,6,7,0,9,10,11,0,13,14,15])
+
+  def test_setitem_single_shard(self):
+    t = self._t(0)
+    t[13] = 99
+    self.assertListEqual(t.tolist(), [0,1,2,3,4,5,6,7,8,9,10,11,12,99,14,15])
+
+  def test_setitem_tensor_value_replicated(self):
+    t = self._t(0)
+    t[2:6] = Tensor([90, 91, 92, 93]).shard(self.device)
+    self.assertListEqual(t.tolist(), [0,1,90,91,92,93,6,7,8,9,10,11,12,13,14,15])
+
+  def test_setitem_tensor_value_sharded_aligned(self):
+    t = self._t(0)
+    t[::4] = Tensor([90, 91, 92, 93]).shard(self.device, axis=0)
+    self.assertListEqual(t.tolist(), [90,1,2,3,91,5,6,7,92,9,10,11,93,13,14,15])
+
 @unittest.skipIf(not_support_multi_device(), "need multi")
 class TestMultiTransformer(unittest.TestCase):
   @needs_second_gpu

test/backend/test_ops.py

@@ -3295,6 +3295,7 @@ class TestOps(unittest.TestCase):
 
 @unittest.skipUnless(is_dtype_supported(dtypes.uchar), f"no uint8 on {Device.DEFAULT}")
 class TestOpsUint8(unittest.TestCase):
+  @unittest.skip("relied on hacks")
   def test_cast(self):
     helper_test_op([(2,3,64,64)], lambda x: x.type(torch.uint8), lambda x: x.cast('uint8'), forward_only=True)
 

test/backend/test_outerworld_call.py

@@ -0,0 +1,19 @@
+import unittest
+from tinygrad import Tensor
+
+class TestOuterCall(unittest.TestCase):
+  def test_outer_call_assign(self):
+    a = Tensor.zeros(10,10).contiguous()
+    b = Tensor.ones(10,10).contiguous()
+    Tensor.realize(a,b)
+
+    pa = a.as_param(0)
+    pb = b.as_param(1)
+    out = Tensor.call(a, b, fxn=pa.assign(pa+pb))
+    out.realize()
+
+    print(a.numpy())
+    assert (a == 1).all().item()
+
+if __name__ == '__main__':
+  unittest.main()

test/backend/test_schedule.py

@@ -1018,7 +1018,8 @@ class TestSchedule(unittest.TestCase):
     a = Tensor.arange(16).contiguous().realize()
     GlobalCounters.reset()
     a[4] = 3
-    # TODO: update when this becomes lazy
+    self.assertEqual(GlobalCounters.kernel_count, 0)
+    a.realize()
     self.assertEqual(GlobalCounters.kernel_count, 1)
     self.assertListEqual(a.tolist(), [0, 1, 2, 3, 3, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15])
 
@@ -1081,6 +1082,14 @@ class TestSchedule(unittest.TestCase):
     new_uop = a.reshape(4,1).realize().uop
     assert new_uop.base.op is Ops.BUFFER
 
+  def test_self_assign_no_empty_kernel(self):
+    for shape in [(3, 3), (4, 4)]:
+      a = Tensor.ones(*shape).contiguous().realize()
+      a.assign(a / 1)
+      run_schedule(check_schedule(a, 0, filter_sink=False))
+      self.assertListEqual(a.tolist(), [[1.]*shape[1]]*shape[0])
+
+class TestLimitBufs(unittest.TestCase):
   @unittest.skipIf(CI and Device.DEFAULT == "NV", "crashes on NV CI")
   def test_limit_bufs_with_var(self):
     N = 31
@@ -1093,12 +1102,16 @@ class TestSchedule(unittest.TestCase):
     for X in range(1,N): root = root + bufs[X][vi] + bufs[X][vj]
     self.assertEqual(root.item(), N * 2)
 
-  def test_self_assign_no_empty_kernel(self):
-    for shape in [(3, 3), (4, 4)]:
-      a = Tensor.ones(*shape).contiguous().realize()
-      a.assign(a / 1)
-      run_schedule(check_schedule(a, 0, filter_sink=False))
-      self.assertListEqual(a.tolist(), [[1.]*shape[1]]*shape[0])
+  def test_limit_bufs_arange_condition(self):
+    # WHERE with arange-based condition (pure index math, no device) and many buffer loads should not crash limit_bufs
+    with Context(MAX_KERNEL_BUFFERS=8):
+      N = 8
+      idx = Tensor.arange(N)
+      base = Tensor.zeros(N)
+      for i in range(4):
+        a, b = Tensor.rand(N).realize(), Tensor.rand(N).realize()
+        base = (idx >= i).where(a + b, base)
+      assert all(x > 0 for x in base.tolist())
 
 class TestSwizzle(unittest.TestCase):
   def test_swizzle_simple(self):

test/backend/test_setitem.py

@@ -36,18 +36,6 @@ class TestSetitem(unittest.TestCase):
     t[:3] *= 10
     self.assertListEqual(t.tolist(), [0, 10, 20, 3, 4, 5, 6, 7, 8, 9])
 
-  def test_setitem_into_unrealized(self):
-    t = Tensor.arange(4).reshape(2, 2)
-    t[1] = 5
-    np.testing.assert_allclose(t.numpy(), [[0, 1], [5, 5]])
-
-  def test_setitem_into_unrealized_sliced_compute(self):
-    # base computation contains SHRINK from prior slicing (like QR decomposition pattern)
-    a = Tensor.arange(6, dtype=dtypes.float).reshape(2, 3)
-    w = a[0] + a[1]  # unrealized ADD with SHRINK in graph: [3, 5, 7]
-    w[1] = 99
-    np.testing.assert_allclose(w.numpy(), [3, 99, 7])
-
   def test_setitem_fancy_on_unrealized_view(self):
     # fancy indexing setitem on unrealized SHRINK view (triggered infinite loop in graph_rewrite)
     base = Tensor.arange(20, dtype=dtypes.float).reshape(4, 5)
@@ -69,10 +57,6 @@ class TestSetitem(unittest.TestCase):
     t = Tensor.zeros(6, dtype=dtypes.float).contiguous().realize()
     with self.assertRaises(RuntimeError): t[2:4] = Tensor([1, 2], dtype=dtypes.int)
 
-  def test_setitem_into_noncontiguous(self):
-    t = Tensor.ones(4)
-    with self.assertRaises(RuntimeError): t[1] = 5
-
   def test_setitem_chained_indexing(self):
     # N[i][j] must work the same as N[i, j]
     N1 = Tensor.zeros((3, 3)).contiguous().realize()
@@ -162,6 +146,8 @@ class TestSetitem(unittest.TestCase):
     @TinyJit
     def f(t:Tensor, a:Tensor):
       t[2:4, 3:5] = a
+      # NOTE: without return t or an explicit realize, it's lazy and not captured
+      return t
 
     for i in range(1, 6):
       t = Tensor.zeros(6, 6).contiguous().realize()

test/external/external_test_hive_reset.py

@@ -0,0 +1,42 @@
+#!/usr/bin/env python3
+import subprocess, sys
+from tinygrad.helpers import getenv
+
+LOOPS = getenv("LOOPS", 10)
+BROKEN = getenv("BROKEN", 0)
+
+BROKEN_KERNEL_SCRIPT = """
+from tinygrad.device import Device
+from tinygrad.runtime.ops_amd import AMDProgram, AMDDevice
+from tinygrad.runtime.support.compiler_amd import compile_hip
+dev = Device["AMD"]
+assert isinstance(dev, AMDDevice) and dev.is_am(), "Need AM driver (not KFD)"
+broken_src = '''
+extern "C" __attribute__((global)) void broken(int* dummy) {
+  volatile int* bad_ptr = (volatile int*)0xDEAD00000000ULL;
+  *bad_ptr = 0x42;
+}
+'''
+broken_lib = compile_hip(broken_src, dev.arch)
+broken_prg = AMDProgram(dev, "broken", broken_lib)
+buf = dev.allocator.alloc(64)
+try:
+  broken_prg(buf, global_size=(1,1,1), local_size=(1,1,1), wait=True)
+  print("  ERROR: Kernel did not fault!")
+except RuntimeError as e:
+  print(f"  Got expected error: {e}")
+"""
+
+for i in range(LOOPS):
+  print(f"=== Running hive_reset.py ({i+1}/{LOOPS}) ===")
+  subprocess.run([sys.executable, "extra/amdpci/hive_reset.py"], check=True)
+  print("=== hive_reset complete ===")
+
+  if BROKEN:
+    print(f"=== Running broken kernel ({i+1}/{LOOPS}) ===")
+    ret = subprocess.run([sys.executable, "-c", BROKEN_KERNEL_SCRIPT])
+    print(f"=== broken kernel exited with code {ret.returncode} ===")
+
+  print(f"=== Running test_tiny.py ({i+1}/{LOOPS}) ===")
+  ret = subprocess.run([sys.executable, "test/test_tiny.py", "TestTiny.test_plus"])
+  print(f"=== test_tiny.py exited with code {ret.returncode} ===")

test/helpers.py

@@ -41,14 +41,18 @@ def assert_jit_cache_len(fxn, expected_len):
     assert type(fxn.jit_cache[0].prg).__name__.endswith('Graph')
     assert len(fxn.jit_cache[0].prg.jit_cache) == expected_len
 
-def rand_for_dtype(dt:DType, size:int):
+def rand_for_dtype(dt:DType, size:int, allow_subnormal=True):
   if dtypes.is_unsigned(dt):
     return np.random.randint(0, 100, size=size, dtype=_to_np_dtype(dt))
   elif dtypes.is_int(dt):
     return np.random.randint(-100, 100, size=size, dtype=_to_np_dtype(dt))
   elif dt == dtypes.bool:
     return np.random.choice([True, False], size=size)
-  return np.random.uniform(-10, 10, size=size).astype(_to_np_dtype(dt))
+  ret = np.random.uniform(-10, 10, size=size).astype(_to_np_dtype(dt))
+  if not allow_subnormal:
+    min_normal = 2.0 ** (2 - (1 << (dtypes.finfo(dt)[0] - 1)))
+    ret = np.where(np.abs(ret) < min_normal, 0, ret)
+  return ret
 
 def timeit(fxn:Callable[..., T], *args, **kwargs) -> tuple[T, float]:
   st = time.perf_counter_ns()

test/mockgpu/amd/README

@@ -4,12 +4,12 @@ Test with `pytest -n12 test/amd/`
 `AMD_LLVM=1 pytest -n12 test/amd/`
 
 * dsl.py -- helpers for the autogen instruction classes in `__init__.py`. should be standalone with init
-* emu.py -- an emulator for RDNA that runs in tinygrad with `AMD=1 MOCKGPU=1 PYTHON_REMU=1`
+* test/mockgpu/amd/emu.py -- an emulator for RDNA that runs in tinygrad with `AMD=1 MOCKGPU=1 PYTHON_REMU=1`
 * generate.py -- extract assembly format + instruction pseudocode from AMD XML + PDF
-* pcode.py -- pseudocode to UOp transformation
+* test/mockgpu/amd/pcode.py -- pseudocode to UOp transformation
 * sqtt.py -- SQTT parser
 
-The code should be as readable and deduplicated as possible. asm and emu shouldn't be required for dsl.
+The code should be as readable and deduplicated as possible. emu (in test/mockgpu/amd/) shouldn't be required for dsl.
 
 The autogen folder is autogenerated from the AMD PDFs with `python3 -m tinygrad.renderer.amd.pdf --arch all`
 

test/mockgpu/amd/emu.py

@@ -67,7 +67,7 @@ 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
+from test.mockgpu.amd.pcode import parse_block, _FUNCS
 
 MASK32 = 0xFFFFFFFF
 

test/mockgpu/helpers.py

@@ -24,7 +24,7 @@ class PythonRemu:
   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
+    from test.mockgpu.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, self.user_data)
 
 def _try_dlopen_remu():

test/null/test_tensor_metadata.py

@@ -1,6 +1,7 @@
 import unittest
 from tinygrad import Tensor, dtypes
 from tinygrad.tensor import _METADATA
+from tinygrad.engine.realize import capturing
 from tinygrad.helpers import Context
 
 class TestTensorMetadata(unittest.TestCase):
@@ -62,6 +63,7 @@ class TestTensorMetadata(unittest.TestCase):
     self.assertEqual(len(si.metadata), 3)
     self.assertEqual(set(m.name for m in si.metadata), {"relu", "sigmoid", "__mul__"})
 
+  @unittest.skip("flaky")
   def test_complex_backward(self):
     x = Tensor.rand(3, requires_grad=True).realize()
     y = Tensor.rand(3, requires_grad=True).realize()
@@ -90,5 +92,25 @@ class TestTensorMetadata(unittest.TestCase):
       si = out.schedule()[-1]
       self.assertEqual(si.metadata, ())
 
+  def _has_metadata(self, h, name):
+    items = []
+    capturing.append(type("", (), {"add": lambda _, ei: items.append(ei)})())
+    try: h.realize()
+    finally: capturing.clear()
+    return any(m.name == name for ei in items for m in ei.metadata)
+
+  def test_metadata_survives_realize_pending_assign(self):
+    shared = Tensor.rand(4)
+    c = Tensor.zeros(8).contiguous().realize()
+    c[:4].assign(shared)
+    self.assertTrue(self._has_metadata(c[:4].relu(), "relu"))
+
+  @unittest.expectedFailure
+  def test_metadata_lost_realize_pending_assign(self):
+    shared = Tensor.rand(4)
+    c = Tensor.zeros(8).contiguous().realize()
+    c[:4].assign(shared)
+    self.assertTrue(self._has_metadata((c[:4] + shared).relu(), "relu"))
+
 if __name__ == '__main__':
   unittest.main()

test/null/test_tqdm.py

@@ -326,7 +326,7 @@ class TestProgressBar(unittest.TestCase):
     for _ in tinytqdm(range(10^7)): pass
     tinytqdm_time = time.perf_counter() - st
 
-    assert tinytqdm_time < 5 * tqdm_time
+    assert tinytqdm_time < 20 * tqdm_time
 
 if __name__ == '__main__':
   unittest.main()

test/null/test_uops_stats.py

@@ -68,28 +68,6 @@ class TestMemoryCount(unittest.TestCase):
     _, mem = get_stats(a.assign(a+a))
     self.assertEqual(mem, 1024*1024*2)  # 1 read + 1 write
 
-  def test_setitem_slice_const(self):
-    t = Tensor.empty(100, dtype=dtypes.int).realize()
-    GlobalCounters.reset()
-    t[20:50] = 3
-    t.realize()
-    self.assertEqual(GlobalCounters.global_mem, 30*4)  # 30 elements written
-
-  def test_setitem_slice_tensor(self):
-    t = Tensor.empty(100, dtype=dtypes.int).realize()
-    v = Tensor.empty(30, dtype=dtypes.int).realize()
-    GlobalCounters.reset()
-    t[20:50] = v
-    t.realize()
-    self.assertEqual(GlobalCounters.global_mem, 30*4*2)  # 30 read + 30 written
-
-  def test_setitem_full(self):
-    t = Tensor.empty(100, dtype=dtypes.int).realize()
-    GlobalCounters.reset()
-    t[:] = 3
-    t.realize()
-    self.assertEqual(GlobalCounters.global_mem, 100*4)  # full buffer written
-
   @unittest.skipIf(Device.DEFAULT == "CPU", "test copy to CPU from other device")
   def test_copyout(self):
     a = Tensor.empty(32, dtype=dtypes.uint8).to("CPU")

test/null/test_viz.py

@@ -364,6 +364,15 @@ def load_profile(lst:list[ProfileEvent]) -> dict:
   return {"dur":total_dur, "peak":global_peak, "layout":layout, "markers":markers}
 
 class TestVizProfiler(BaseTestViz):
+  def test_transfer_uses_copy_device(self):
+    a = Tensor.ones(1, device="NULL").contiguous().realize()
+    a.to("NULL:1").realize()
+    range_events = [e for e in cpu_events if isinstance(e, ProfileRangeEvent)]
+    compute_events = [e for e in range_events if e.device == "NULL"]
+    copy_events = [e for e in range_events if e.device.endswith(":COPY")]
+    self.assertGreater(len(compute_events), 0, "expected compute events on base device")
+    self.assertGreater(len(copy_events), 0, "transfer must produce events with ':COPY' device suffix")
+
   def test_node(self):
     prof = [ProfileRangeEvent(device='NV', name='E_2', st=decimal.Decimal(1000), en=decimal.Decimal(1010)),
             ProfileDeviceEvent(device='NV', tdiff=decimal.Decimal(-1000))]
@@ -574,6 +583,7 @@ class TestVizMemoryLayout(BaseTestViz):
     user_cnt = [len(b["arg"]["users"]) for b in buffers if b["arg"].get("users")]
     self.assertEqual(len(user_cnt), len(programs))
 
+  @unittest.skip("flaky")
   def test_inflight_buf(self):
     a = Tensor.empty(1, device="NULL")
     n = 4

test/testextra/test_tk.py

@@ -5,21 +5,18 @@ from tinygrad.uop.ops import UOp, Ops
 from tinygrad.engine.realize import get_runner
 from tinygrad.engine.schedule import ExecItem
 from tinygrad.engine.jit import TinyJit
-from tinygrad.helpers import CI
 import numpy as np
 
 from extra.thunder.tiny.tk import WARP_THREADS
 from extra.thunder.tiny.tk.kernel import Kernel
 from extra.thunder.tiny.tk.tiles import ST_16X32, RT_16X32, RT_16X16, TileLayout
 
-@unittest.skipIf(CI or Device.DEFAULT not in ["AMD"], "only amd")
 class TestTK(unittest.TestCase):
   def setUp(self):
-    arch = Device["AMD"].arch
+    arch = getattr(Device[Device.DEFAULT].renderer, "arch", "")
     if not arch.startswith("gfx9"):
       self.skipTest(f"arch {arch} not supported")
 
-  @unittest.skipIf(CI, "no wmma in ci")
   def test_simple_matmul(self):
     N = 8192
     BLOCK_SIZE = 64
@@ -73,7 +70,6 @@ class TestTK(unittest.TestCase):
 
     np.testing.assert_allclose(c.numpy(), ref.numpy())
 
-  @unittest.skipIf(CI, "no wmma in ci")
   def test_simple_matmul_transposed(self):
     N = 8192
     BLOCK_N, BLOCK_M, BLOCK_K = 64, 64, 128

test/unit/test_assign.py

@@ -756,6 +756,26 @@ class TestAssignOrdering(unittest.TestCase):
     self.assertEqual(buf[0:1, :].sum().item(), 4)
     self.assertEqual(buf[1:2, :].sum().item(), 8)
 
+  def test_multi_step_assign_read_write_same_buffer(self):
+    """Assign to m and param reading b, then update b, across multiple steps.
+    This is the optimizer bias-correction pattern from issue #13600: m accumulates,
+    param is updated using m/(1-b), and b is updated via *= after the reads."""
+    b = Tensor([0.5]).contiguous().realize()
+    m = Tensor([0.0]).contiguous().realize()
+    param = Tensor([1.0]).contiguous().realize()
+    for _ in range(10):
+      m.assign(0.9 * m + 0.1)
+      param.assign(param - m / (1 - b))
+      b *= 0.9
+      Tensor.realize(param, m, b)
+    # numpy reference
+    b_np, m_np, p_np = 0.5, 0.0, 1.0
+    for _ in range(10):
+      m_np = 0.9 * m_np + 0.1
+      p_np = p_np - m_np / (1 - b_np)
+      b_np *= 0.9
+    np.testing.assert_allclose(param.item(), p_np, atol=1e-5)
+
   def test_multiple_slice_assigns_then_read(self):
     """Multiple non-overlapping slice assigns then read."""
     buf = Tensor.zeros(4).contiguous().realize()

test/unit/test_disk_tensor.py

@@ -456,6 +456,7 @@ class TestDiskTensor(TempDirTestCase):
     np.testing.assert_equal(t1.numpy(), np.arange(128, dtype=np.uint8))
     np.testing.assert_equal(t2.numpy(), np.arange(64, dtype=np.uint8))
 
+  @unittest.skip("fails with setup_python_cap run")
   def test_disk_open_failure_state(self):
     from tinygrad.runtime.ops_disk import DiskDevice
     fn = pathlib.Path(self.tmp("dt_open_failure"))
@@ -476,6 +477,7 @@ class TestDiskTensor(TempDirTestCase):
     t2.to("CPU").realize()
     assert disk_device.size == 200
 
+  @unittest.skip("fails with setup_python_cap run")
   def test_disk_permission_error(self):
     fn = pathlib.Path(self.tmp("dt_permission"))
     fn.write_bytes(bytes(range(256)))

test/unit/test_indexing.py

@@ -1000,7 +1000,7 @@ def assert_backward_eq(tensor: Tensor, indexer):
 def get_set_tensor(indexed: Tensor, indexer):
   set_size = indexed[indexer].shape
   set_count = indexed[indexer].numel()
-  set_tensor = Tensor.randint(set_count, high=set_count).reshape(set_size) #.cast(dtypes.float64)
+  set_tensor = Tensor.randint(set_count, high=set_count).reshape(set_size).cast(indexed.dtype)
   return set_tensor
 
 @slow

test/unit/test_setitem_schedule.py

@@ -0,0 +1,150 @@
+import unittest
+from tinygrad import Tensor, dtypes, GlobalCounters
+
+class TestSetitemInto(unittest.TestCase):
+  def test_setitem_into_unrealized(self):
+    GlobalCounters.reset()
+    t = Tensor.arange(4, dtype=dtypes.int32).reshape(2, 2)
+    self.assertEqual(GlobalCounters.kernel_count, 0)
+    t[1] = 5
+    self.assertEqual(GlobalCounters.kernel_count, 0)
+    t.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertEqual(GlobalCounters.global_mem, 16)
+    t[1].realize()
+    t.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertListEqual(t.tolist(), [[0, 1], [5, 5]])
+
+  def test_setitem_into_unrealized_sliced_compute(self):
+    # base computation contains SHRINK from prior slicing (like QR decomposition pattern)
+    GlobalCounters.reset()
+    a = Tensor.arange(8, dtype=dtypes.int32).reshape(2, 4)
+    w = a[0] + a[1]  # unrealized ADD with SHRINK in graph: [4, 6, 8, 10]
+    self.assertEqual(GlobalCounters.kernel_count, 0)
+    w[1] = 99
+    self.assertEqual(GlobalCounters.kernel_count, 0)
+    w.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertEqual(GlobalCounters.global_mem, 4*4)
+    self.assertListEqual(w.tolist(), [4, 99, 8, 10])
+
+  def test_setitem_into_empty(self):
+    GlobalCounters.reset()
+    t = Tensor.empty(4, dtype=dtypes.int32)
+    t[1] = 5
+    self.assertEqual(GlobalCounters.kernel_count, 0)
+    t.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    # TODO: this can be just 4 if empty goes through is_realized setitem path
+    self.assertEqual(GlobalCounters.global_mem, 4*(3*2+1)) # 3 elements had +1, 1 is assigned directly
+    t[1].realize()
+    t.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertEqual(t[1].item(), 5)
+
+  def test_setitem_into_empty_alu(self):
+    GlobalCounters.reset()
+    t = Tensor.empty(4, dtype=dtypes.int32) + 1
+    self.assertEqual(GlobalCounters.kernel_count, 0)
+    t[1] = 5
+    self.assertEqual(GlobalCounters.kernel_count, 0)
+    t.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertEqual(GlobalCounters.global_mem, 4*(3*2+1)) # 3 elements had +1, 1 is assigned directly
+    t[1].realize()
+    t.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertEqual(t[1].item(), 5)
+
+  def test_setitem_into_tensor(self):
+    t = Tensor([1, 2, 3, 4], dtype=dtypes.int32).realize()
+    GlobalCounters.reset()
+    t[1] = 5
+    self.assertEqual(GlobalCounters.kernel_count, 0)
+    t[1].realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertEqual(GlobalCounters.global_mem, 4)
+    t.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertListEqual(t.tolist(), [1, 5, 3, 4])
+
+  def test_setitem_into_tensor_alu(self):
+    t = Tensor([1, 2, 3, 4], dtype=dtypes.int32).realize() + 1
+    GlobalCounters.reset()
+    t[1] = 5
+    self.assertEqual(GlobalCounters.kernel_count, 0)
+    t[1].realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertEqual(GlobalCounters.global_mem, 4*(3*2+1)) # 3 elements had +1, 1 is assigned directly
+    t[1].realize()
+    t.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertListEqual(t.tolist(), [2, 5, 4, 5])
+
+  def test_setitem_into_cont(self):
+    GlobalCounters.reset()
+    t = Tensor.ones(4, dtype=dtypes.int32)
+    t[1] = 5
+    self.assertEqual(GlobalCounters.kernel_count, 0)
+    t.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertEqual(GlobalCounters.global_mem, 4*4)
+    t[1].realize()
+    t.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertListEqual(t.tolist(), [1, 5, 1, 1])
+
+  def test_setitem_into_const_alu(self):
+    GlobalCounters.reset()
+    t = Tensor.ones(4, dtype=dtypes.int32) + 1
+    t[1] = 5
+    self.assertEqual(GlobalCounters.kernel_count, 0)
+    t.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertEqual(GlobalCounters.global_mem, 4*4)
+    t[1].realize()
+    t.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertListEqual(t.tolist(), [2, 5, 2, 2])
+
+  def test_setitem_into_arange(self):
+    # NOTE: arange has no real buffer, but assigning to it is fine
+    GlobalCounters.reset()
+    t = Tensor.arange(4, dtype=dtypes.int32)
+    t[1] = 5
+    self.assertEqual(GlobalCounters.kernel_count, 0)
+    t.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertListEqual(t.tolist(), [0, 5, 2, 3])
+
+  def test_setitem_slice_const(self):
+    t = Tensor.zeros(100, dtype=dtypes.int32).contiguous().realize()
+    GlobalCounters.reset()
+    t[20:50] = 3
+    self.assertEqual(GlobalCounters.kernel_count, 0)
+    t.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertEqual(GlobalCounters.global_mem, 30*4)  # 30 elements written
+
+  def test_setitem_slice_tensor(self):
+    t = Tensor.zeros(100, dtype=dtypes.int32).contiguous().realize()
+    v = Tensor.zeros(30, dtype=dtypes.int32).contiguous().realize()
+    GlobalCounters.reset()
+    t[20:50] = v
+    self.assertEqual(GlobalCounters.kernel_count, 0)
+    t.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertEqual(GlobalCounters.global_mem, 30*4*2)  # 30 read + 30 written
+
+  def test_setitem_full(self):
+    t = Tensor.zeros(100, dtype=dtypes.int32).contiguous().realize()
+    GlobalCounters.reset()
+    t[:] = 3
+    self.assertEqual(GlobalCounters.kernel_count, 0)
+    t.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertEqual(GlobalCounters.global_mem, 100*4)  # full buffer written
+
+if __name__ == '__main__':
+  unittest.main()

tinygrad/engine/schedule.py

@@ -29,7 +29,9 @@ def create_schedule(sched_sink:UOp) -> tuple[list[ExecItem], UOp]:
       assert k.op in {Ops.CALL, Ops.END}, f"AFTER src[1] should be KERNEL or END, not {k.op}"
       in_degree.setdefault(k, 0)
       if k.op is Ops.END: assert k.src[0].op is Ops.CALL, f"END src[0] should be KERNEL, not {k.src[0].op}"
-      for s in k.src[0].src[1:] if k.op is Ops.END else k.src[1:]:
+      # WAR deps from rangeify are stored in AFTER src[2:]
+      kernel_deps = k.src[0].src[1:] if k.op is Ops.END else k.src[1:]
+      for s in kernel_deps + u.src[2:]:
         match (s := _unwrap_src(s)).op:
           case Ops.AFTER:
             children.setdefault(s.src[1], []).append(k)

tinygrad/helpers.py

@@ -182,6 +182,7 @@ CACHELEVEL, IGNORE_BEAM_CACHE, DEVECTORIZE = ContextVar("CACHELEVEL", 2), Contex
 VALIDATE_WITH_CPU, DISABLE_FAST_IDIV = ContextVar("VALIDATE_WITH_CPU", 0), ContextVar("DISABLE_FAST_IDIV", 0)
 CORRECT_DIVMOD_FOLDING, FUSE_OPTIM = ContextVar("CORRECT_DIVMOD_FOLDING", 0), ContextVar("FUSE_OPTIM", 0)
 ALLOW_DEVICE_USAGE, MAX_BUFFER_SIZE = ContextVar("ALLOW_DEVICE_USAGE", 1), ContextVar("MAX_BUFFER_SIZE", 0)
+MAX_KERNEL_BUFFERS = ContextVar("MAX_KERNEL_BUFFERS", 0)
 EMULATE, EMULATED_DTYPES = ContextVar("EMULATE", ""), ContextVar("EMULATED_DTYPES", "")
 CPU_COUNT = ContextVar("CPU_COUNT", max(1, len(os.sched_getaffinity(0)) if hasattr(os, "sched_getaffinity") else (os.cpu_count() or 1)))
 # Compilers
@@ -189,7 +190,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, AMD_ISEL, AMD_ASM = ContextVar("AMD_CC", ""), ContextVar("AMD_LLVM", 0), ContextVar("AMD_HIPCC", 0), ContextVar("AMD_ISEL", 0), ContextVar("AMD_ASM", 0)
+AMD_CC, AMD_LLVM, AMD_HIPCC  = ContextVar("AMD_CC", ""), ContextVar("AMD_LLVM", 0), ContextVar("AMD_HIPCC", 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,9 +444,6 @@ 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/elf.py

@@ -9,98 +9,11 @@ from tinygrad.renderer.amd.dsl import Reg, FixedBitField
 from tinygrad.runtime.autogen.amd.rdna3.ins import s_code_end # same encoding as RDNA4
 from tinygrad.runtime.autogen.amd.cdna.ins import s_nop as s_nop_cdna
 
-def put(dst:bytearray, off:int, data:bytes) -> None:
-  end = off + len(data)
-  if end > len(dst): raise ValueError("write past end of buffer")
-  dst[off:end] = data
-
-def create_elf(prg:bytes, kd:dict, arch:str) -> bytes:
-  is_cdna, is_rdna4 = arch == "cdna", arch == "rdna4"
-  padding_inst = (s_nop_cdna(0) if is_cdna else s_code_end()).to_bytes()
-  text = prg + padding_inst * ((hsa.AMD_ISA_ALIGN_BYTES - len(prg) % hsa.AMD_ISA_ALIGN_BYTES) % hsa.AMD_ISA_ALIGN_BYTES)
-  text_offset = round_up(ctypes.sizeof(libc.Elf64_Ehdr), hsa.AMD_ISA_ALIGN_BYTES)
-  rodata_offset = text_offset + len(text)
-
-  # ** pack rodata object
-  desc = amdgpu_kd.llvm_amdhsa_kernel_descriptor_t()
-  desc.group_segment_fixed_size = kd.get("group_segment_fixed_size", 0)
-  desc.private_segment_fixed_size = kd.get("private_segment_fixed_size", 0)
-  desc.kernarg_size = kd.get("kernarg_size", 0)
-  desc.kernel_code_entry_byte_offset = text_offset-rodata_offset
-  # rsrc1
-  vgpr_granule = max(0, (kd["next_free_vgpr"] + 7) // 8 - 1)
-  # CDNA: add 6 for VCC(2) + FLAT_SCRATCH(2) + XNACK_MASK(2)
-  # next_free_sgpr is unused in RDNA
-  # NOTE: CU mode is the default, it seems faster and simpler
-  sgpr_granule = max(0, ceildiv(kd["next_free_sgpr"] + 6, 8) - 1) if is_cdna else 0
-  desc.compute_pgm_rsrc1 = (vgpr_granule << amdgpu_kd.COMPUTE_PGM_RSRC1_GRANULATED_WORKITEM_VGPR_COUNT_SHIFT |
-                            sgpr_granule << amdgpu_kd.COMPUTE_PGM_RSRC1_GRANULATED_WAVEFRONT_SGPR_COUNT_SHIFT |
-                            kd.get("float_round_mode_32", 0) << amdgpu_kd.COMPUTE_PGM_RSRC1_FLOAT_ROUND_MODE_32_SHIFT |
-                            kd.get("float_round_mode_16_64", 0) << amdgpu_kd.COMPUTE_PGM_RSRC1_FLOAT_ROUND_MODE_16_64_SHIFT |
-                            kd.get("float_denorm_mode_32", 0) << amdgpu_kd.COMPUTE_PGM_RSRC1_FLOAT_DENORM_MODE_32_SHIFT |
-                            kd.get("float_denorm_mode_16_64", 3) << amdgpu_kd.COMPUTE_PGM_RSRC1_FLOAT_DENORM_MODE_16_64_SHIFT |
-                            kd.get("dx10_clamp", 0 if is_rdna4 else 1) << amdgpu_kd.COMPUTE_PGM_RSRC1_GFX6_GFX11_ENABLE_DX10_CLAMP_SHIFT |
-                            kd.get("ieee_mode", 0 if is_rdna4 else 1) << amdgpu_kd.COMPUTE_PGM_RSRC1_GFX6_GFX11_ENABLE_IEEE_MODE_SHIFT |
-                            kd.get("fp16_overflow", 0) << amdgpu_kd.COMPUTE_PGM_RSRC1_GFX9_PLUS_FP16_OVFL_SHIFT |
-                            (0 if is_cdna else kd.get("workgroup_processor_mode", 0)) << amdgpu_kd.COMPUTE_PGM_RSRC1_GFX10_PLUS_WGP_MODE_SHIFT |
-                            (0 if is_cdna else kd.get("memory_ordered", 1)) << amdgpu_kd.COMPUTE_PGM_RSRC1_GFX10_PLUS_MEM_ORDERED_SHIFT |
-                            (0 if is_cdna else kd.get("forward_progress", 0)) << amdgpu_kd.COMPUTE_PGM_RSRC1_GFX10_PLUS_FWD_PROGRESS_SHIFT)
-  # rsrc2
-  desc.compute_pgm_rsrc2 = (kd.get("enable_private_segment", 0) << amdgpu_kd.COMPUTE_PGM_RSRC2_ENABLE_PRIVATE_SEGMENT_SHIFT |
-                            kd.get("user_sgpr_count", 0) << amdgpu_kd.COMPUTE_PGM_RSRC2_USER_SGPR_COUNT_SHIFT |
-                            kd.get("system_sgpr_workgroup_id_x", 1) << amdgpu_kd.COMPUTE_PGM_RSRC2_ENABLE_SGPR_WORKGROUP_ID_X_SHIFT |
-                            kd.get("system_sgpr_workgroup_id_y", 0) << amdgpu_kd.COMPUTE_PGM_RSRC2_ENABLE_SGPR_WORKGROUP_ID_Y_SHIFT |
-                            kd.get("system_sgpr_workgroup_id_z", 0) << amdgpu_kd.COMPUTE_PGM_RSRC2_ENABLE_SGPR_WORKGROUP_ID_Z_SHIFT |
-                            kd.get("system_sgpr_workgroup_info", 0) << amdgpu_kd.COMPUTE_PGM_RSRC2_ENABLE_SGPR_WORKGROUP_INFO_SHIFT |
-                            kd.get("system_vgpr_workitem_id", 0) << amdgpu_kd.COMPUTE_PGM_RSRC2_ENABLE_VGPR_WORKITEM_ID_SHIFT)
-  # rsrc3
-  if is_cdna:
-    amdhsa_accum_offset = ((kd.get("accum_offset", 4) // 4) - 1) & amdgpu_kd.COMPUTE_PGM_RSRC3_GFX90A_ACCUM_OFFSET
-    desc.compute_pgm_rsrc3 = amdhsa_accum_offset << amdgpu_kd.COMPUTE_PGM_RSRC3_GFX90A_ACCUM_OFFSET_SHIFT
-  else:
-    desc.compute_pgm_rsrc3 = kd.get("shared_vgpr_count", 0) << amdgpu_kd.COMPUTE_PGM_RSRC3_GFX10_GFX11_SHARED_VGPR_COUNT_SHIFT
-  # kernel code properties
-  desc.kernel_code_properties = (kd.get("user_sgpr_dispatch_ptr", 0) << amdgpu_kd.KERNEL_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR_SHIFT |
-                                 kd.get("user_sgpr_queue_ptr", 0) << amdgpu_kd.KERNEL_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR_SHIFT |
-                                 kd.get("user_sgpr_kernarg_segment_ptr", 0) << amdgpu_kd.KERNEL_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR_SHIFT |
-                                 kd.get("user_sgpr_dispatch_id", 0) << amdgpu_kd.KERNEL_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID_SHIFT |
-                                 kd.get("user_sgpr_private_segment_size",0) << amdgpu_kd.KERNEL_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE_SHIFT |
-                                 kd.get("wavefront_size32", 0 if is_cdna else 1) << amdgpu_kd.KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32_SHIFT |
-                                 kd.get("uses_dynamic_stack", 0) << amdgpu_kd.KERNEL_CODE_PROPERTY_USES_DYNAMIC_STACK_SHIFT)
-  rodata = bytes(desc)
-
-  # ** pack elf sections
-  sh_names:list[int] = []
-  strtab = bytearray(b"\x00")
-  for name in [".text", ".rodata", ".strtab"]:
-    sh_names.append(len(strtab))
-    strtab += name.encode("ascii") + b"\x00"
-
-  rodata_offset = round_up(text_offset+(text_size:=len(text)), hsa.AMD_KERNEL_CODE_ALIGN_BYTES)
-  strtab_offset = rodata_offset+(rodata_size:=len(rodata))
-  shdr_offset   = strtab_offset+(strtab_size:=len(strtab))
-
-  sections = [(libc.SHT_PROGBITS, libc.SHF_ALLOC | libc.SHF_EXECINSTR, text_offset, text_offset, text_size),
-              (libc.SHT_PROGBITS, libc.SHF_ALLOC, rodata_offset, rodata_offset, rodata_size),
-              (libc.SHT_STRTAB, 0, 0, strtab_offset, strtab_size)]
-  shdrs = (libc.Elf64_Shdr * len(sections))()
-  for i,s in enumerate(sections): shdrs[i] = libc.Elf64_Shdr(sh_names[i], *s)
-
-  ehdr = libc.Elf64_Ehdr()
-  ehdr.e_shoff, ehdr.e_shnum, ehdr.e_shstrndx = shdr_offset, len(sections), 2
-
-  elf = bytearray(shdr_offset + ctypes.sizeof(shdrs))
-  put(elf, 0, bytes(ehdr))
-  put(elf, text_offset, text)
-  put(elf, rodata_offset, rodata)
-  put(elf, strtab_offset, strtab)
-  put(elf, shdr_offset, bytes(shdrs))
-  return bytes(elf)
-
 _arch_map = {"gfx9": "cdna", "gfx10": "rdna3", "gfx11": "rdna3", "gfx12": "rdna4"}
 def do_assemble_amd(ctx, prg:UOp, lin:UOp) -> UOp:
   insts = [u.arg for u in lin.src]
-  # scan for max vgpr/sgpr
+
+  # ** scan for max vgpr/sgpr
   max_vgpr, max_sgpr = 0, 0
   for inst in insts:
     for name, field in inst._fields:
@@ -109,7 +22,8 @@ def do_assemble_amd(ctx, prg:UOp, lin:UOp) -> UOp:
       if not isinstance(val, Reg): continue
       if 256 <= val.offset < 512: max_vgpr = max(max_vgpr, (val.offset - 256) + val.sz)
       elif val.offset < 106: max_sgpr = max(max_sgpr, val.offset + val.sz)
-  # scan sink for metadata
+
+  # ** scan sink for metadata
   sink, n_bufs, n_vars, lds_size, gids = prg.src[0], 0, 0, 0, set()
   for u in sink.toposort():
     if u.op is Ops.PARAM: n_bufs += 1
@@ -119,9 +33,65 @@ def do_assemble_amd(ctx, prg:UOp, lin:UOp) -> UOp:
   src = "\n".join(str(inst) for inst in insts)
   code_bytes = b"".join(inst.to_bytes() for inst in insts)
   arch = next(v for k, v in _arch_map.items() if ctx.arch.startswith(k))
-  kd = {"kernarg_size":n_bufs*8+n_vars*4, "group_segment_fixed_size":lds_size,
-        "user_sgpr_kernarg_segment_ptr":1, "user_sgpr_count":2,
-        "system_sgpr_workgroup_id_x":int(0 in gids), "system_sgpr_workgroup_id_y":int(1 in gids), "system_sgpr_workgroup_id_z":int(2 in gids),
-        "next_free_vgpr":round_up(max_vgpr, 8), "next_free_sgpr":round_up(max_sgpr, 8)}
-  binary = create_elf(code_bytes, kd, arch)
+  is_cdna, is_rdna4 = arch == "cdna", arch == "rdna4"
+
+  # ** pad text to ISA alignment
+  padding_inst = (s_nop_cdna(0) if is_cdna else s_code_end()).to_bytes()
+  text = code_bytes + padding_inst * ((hsa.AMD_ISA_ALIGN_BYTES - len(code_bytes) % hsa.AMD_ISA_ALIGN_BYTES) % hsa.AMD_ISA_ALIGN_BYTES)
+  text_offset = round_up(ctypes.sizeof(libc.Elf64_Ehdr), hsa.AMD_ISA_ALIGN_BYTES)
+
+  # ** pack kernel descriptor (rodata)
+  next_free_vgpr, next_free_sgpr = round_up(max_vgpr, 8), round_up(max_sgpr, 8)
+  vgpr_granule = max(0, (next_free_vgpr + 7) // 8 - 1)
+  # CDNA: add 6 for VCC(2) + FLAT_SCRATCH(2) + XNACK_MASK(2), next_free_sgpr is unused in RDNA.
+  sgpr_granule = max(0, ceildiv(next_free_sgpr + 6, 8) - 1) if is_cdna else 0
+  desc = amdgpu_kd.llvm_amdhsa_kernel_descriptor_t()
+  desc.group_segment_fixed_size = lds_size
+  desc.kernarg_size = n_bufs * 8 + n_vars * 4
+  desc.kernel_code_entry_byte_offset = -len(text)
+
+  # https://llvm.org/docs/AMDGPUUsage.html#amdgpu-amdhsa-compute-pgm-rsrc1-gfx6-gfx12-table
+  # NOTE: CU mode is the default
+  desc.compute_pgm_rsrc1 = (vgpr_granule << amdgpu_kd.COMPUTE_PGM_RSRC1_GRANULATED_WORKITEM_VGPR_COUNT_SHIFT |
+                            sgpr_granule << amdgpu_kd.COMPUTE_PGM_RSRC1_GRANULATED_WAVEFRONT_SGPR_COUNT_SHIFT |
+                            3 << amdgpu_kd.COMPUTE_PGM_RSRC1_FLOAT_DENORM_MODE_16_64_SHIFT |
+                            (0 if is_rdna4 else 1) << amdgpu_kd.COMPUTE_PGM_RSRC1_GFX6_GFX11_ENABLE_DX10_CLAMP_SHIFT |
+                            (0 if is_rdna4 else 1) << amdgpu_kd.COMPUTE_PGM_RSRC1_GFX6_GFX11_ENABLE_IEEE_MODE_SHIFT |
+                            (0 if is_cdna else 1) << amdgpu_kd.COMPUTE_PGM_RSRC1_GFX10_PLUS_MEM_ORDERED_SHIFT)
+  desc.compute_pgm_rsrc2 = (2 << amdgpu_kd.COMPUTE_PGM_RSRC2_USER_SGPR_COUNT_SHIFT |
+                            int(0 in gids) << amdgpu_kd.COMPUTE_PGM_RSRC2_ENABLE_SGPR_WORKGROUP_ID_X_SHIFT |
+                            int(1 in gids) << amdgpu_kd.COMPUTE_PGM_RSRC2_ENABLE_SGPR_WORKGROUP_ID_Y_SHIFT |
+                            int(2 in gids) << amdgpu_kd.COMPUTE_PGM_RSRC2_ENABLE_SGPR_WORKGROUP_ID_Z_SHIFT)
+  desc.kernel_code_properties = (1 << amdgpu_kd.KERNEL_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR_SHIFT |
+                                 (0 if is_cdna else 1) << amdgpu_kd.KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32_SHIFT)
+  rodata = bytes(desc)
+
+  # ** pack ELF
+  sh_names:list[int] = []
+  strtab = bytearray(b"\x00")
+  for name in [".text", ".rodata", ".strtab"]:
+    sh_names.append(len(strtab))
+    strtab += name.encode("ascii") + b"\x00"
+
+  rodata_offset = round_up(text_offset + (text_size := len(text)), hsa.AMD_KERNEL_CODE_ALIGN_BYTES)
+  strtab_offset = rodata_offset + (rodata_size := len(rodata))
+  shdr_offset   = strtab_offset + (strtab_size := len(strtab))
+
+  sections = [(libc.SHT_PROGBITS, libc.SHF_ALLOC | libc.SHF_EXECINSTR, text_offset, text_offset, text_size),
+              (libc.SHT_PROGBITS, libc.SHF_ALLOC, rodata_offset, rodata_offset, rodata_size),
+              (libc.SHT_STRTAB, 0, 0, strtab_offset, strtab_size)]
+  shdrs = (libc.Elf64_Shdr * len(sections))()
+  for i, s in enumerate(sections): shdrs[i] = libc.Elf64_Shdr(sh_names[i], *s)
+
+  ehdr = libc.Elf64_Ehdr()
+  ehdr.e_shoff, ehdr.e_shnum, ehdr.e_shstrndx = shdr_offset, len(sections), 2
+
+  elf = bytearray(shdr_offset + ctypes.sizeof(shdrs))
+  elf[0:ctypes.sizeof(ehdr)] = bytes(ehdr)
+  elf[text_offset:text_offset+text_size] = text
+  elf[rodata_offset:rodata_offset+rodata_size] = rodata
+  elf[strtab_offset:strtab_offset+strtab_size] = strtab
+  elf[shdr_offset:shdr_offset+ctypes.sizeof(shdrs)] = bytes(shdrs)
+  binary = bytes(elf)
+
   return prg.replace(src=prg.src[:3]+(UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=binary)))

tinygrad/renderer/amd/isel.py

@@ -1,325 +0,0 @@
-# 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/renderer.py

@@ -1,427 +0,0 @@
-# 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/runtime/ops_amd.py

@@ -8,11 +8,9 @@ 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, AMD_ISEL, AMD_ASM, ceildiv, unwrap
+from tinygrad.helpers import VIZ, AMD_CC, AMD_LLVM, AMD_HIPCC, 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
@@ -972,9 +970,7 @@ 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),
-                             (functools.partial(AMDISELRenderer, self.arch), AMD_ISEL),
-                             (functools.partial(AMDAssemblyRenderer, self.arch), AMD_ASM)], ctrl_var=AMD_CC)
+                             (functools.partial(AMDHIPCCRenderer, self.arch), AMD_HIPCC)], 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/ops_metal.py

@@ -177,7 +177,7 @@ class MetalAllocator(LRUAllocator[MetalDevice]):
     # There is no real metal multidevice support for now, so transfer is used only for tests.
     src_dev.synchronize()
   def _cp_mv(self, dst, src, prof_desc):
-    with cpu_profile(prof_desc, self.dev.device): dst[:] = src
+    with cpu_profile(prof_desc, f"{self.dev.device}:COPY"): dst[:] = src
   def _as_buffer(self, src:MetalBuffer) -> memoryview:
     self.dev.synchronize()
     return to_mv(src.buf.contents(), src.size + src.offset)[src.offset:]

tinygrad/runtime/ops_null.py

@@ -24,7 +24,7 @@ class NullAllocator(Allocator['NullDevice']):
   def _copyout(self, dest:memoryview, src):
     if not NULL_ALLOW_COPYOUT: raise RuntimeError("no copyout on NULL")
   def _transfer(self, dest, src, sz:int, src_dev, dest_dev):
-    with cpu_profile(f"{src_dev.device} -> {dest_dev.device}", self.dev.device): pass
+    with cpu_profile(f"{src_dev.device} -> {dest_dev.device}", f"{self.dev.device}:COPY"): pass
   def _offset(self, buf, offset:int, size:int): pass
 
 class NullGraph(MultiGraphRunner):

tinygrad/runtime/ops_nv.py

@@ -753,7 +753,7 @@ class NVDevice(HCQCompiled[NVSignal]):
     self.iface.rm_control(self.profiler, nv_gpu.NVB0CC_CTRL_CMD_POWER_REQUEST_FEATURES, power_params)
 
     self.pma_buf = self.iface.alloc(getenv("PMA_BUFFER_SIZE", 512) << 20, uncached=True, cpu_cached=True, cpu_access=True)
-    self.pma_bytes = self.iface.alloc(0x1000, uncached=True, cpu_cached=True, read_only=True)
+    self.pma_bytes = self.iface.alloc(0x1000, uncached=True, cpu_cached=True, cpu_access=True, read_only=True)
     self.pma_rptr = 0
 
     pma_stream = nv_gpu.struct_NVB0CC_CTRL_ALLOC_PMA_STREAM_PARAMS(hMemPmaBuffer=self.pma_buf.meta.hMemory,

tinygrad/runtime/ops_qcom.py

@@ -329,7 +329,7 @@ class QCOMAllocator(HCQAllocatorBase):
     return self.dev._gpu_map(opts.external_ptr, size, image=opts.image) if opts.external_ptr else self.dev._gpu_alloc(size, image=opts.image)
 
   def _do_copy(self, src_addr, dest_addr, src_size, real_size, src_stride, dest_stride, prof_text, dest_off=0, src_off=0):
-    with cpu_profile(prof_text, self.dev.device):
+    with cpu_profile(prof_text, f"{self.dev.device}:COPY"):
       while src_off < src_size:
         ctypes.memmove(dest_addr+dest_off, src_addr+src_off, real_size)
         src_off, dest_off = src_off+src_stride, dest_off+dest_stride

tinygrad/runtime/support/hcq.py

@@ -516,7 +516,7 @@ class HCQAllocator(HCQAllocatorBase, Generic[HCQDeviceType]):
   def _copyin(self, dest:HCQBuffer, src:memoryview):
     if self.dev.hw_copy_queue_t is None:
       self.dev.synchronize()
-      with cpu_profile(f'TINY -> {self.dev.device}', self.dev.device): ctypes.memmove(int(dest.va_addr), from_mv(src), len(src))
+      with cpu_profile(f'TINY -> {self.dev.device}', f"{self.dev.device}:COPY"): ctypes.memmove(int(dest.va_addr), from_mv(src), len(src))
       return
 
     with hcq_profile(self.dev, queue_type=self.dev.hw_copy_queue_t, desc=f"TINY -> {self.dev.device}", enabled=PROFILE, dev_suff="SDMA:0"):
@@ -550,7 +550,7 @@ class HCQAllocator(HCQAllocatorBase, Generic[HCQDeviceType]):
   def _copyout(self, dest:memoryview, src:HCQBuffer):
     self.dev.synchronize()
     if self.dev.hw_copy_queue_t is None:
-      with cpu_profile(f'{self.dev.device} -> TINY', self.dev.device): ctypes.memmove(from_mv(dest), int(src.va_addr), len(dest))
+      with cpu_profile(f'{self.dev.device} -> TINY', f"{self.dev.device}:COPY"): ctypes.memmove(from_mv(dest), int(src.va_addr), len(dest))
       return
 
     with hcq_profile(self.dev, queue_type=self.dev.hw_copy_queue_t, desc=f"{self.dev.device} -> TINY", enabled=PROFILE, dev_suff="SDMA:0"):

tinygrad/schedule/rangeify.py

@@ -4,7 +4,7 @@ from tinygrad.dtype import dtypes, PtrDType, ImageDType, AddrSpace
 from tinygrad.uop.ops import PatternMatcher, UPat, Ops, UOp, resolve, GroupOp, _substitute, KernelInfo, pm_gate_kernel_sink
 from tinygrad.uop.ops import graph_rewrite, identity_element, sint, AxisType, BottomUpGate, _remove_all_tags, range_str
 from tinygrad.uop.symbolic import symbolic
-from tinygrad.helpers import argsort, prod, all_same, getenv, flatten, dedup, all_int, DEBUG, SPLIT_REDUCEOP, DEBUG_RANGEIFY, VIZ
+from tinygrad.helpers import argsort, prod, all_same, getenv, flatten, dedup, all_int, DEBUG, SPLIT_REDUCEOP, DEBUG_RANGEIFY, VIZ, MAX_KERNEL_BUFFERS
 from tinygrad.helpers import PCONTIG, partition, get_single_element
 from tinygrad.codegen.simplify import pm_flatten_range, pm_reduce_simplify
 from tinygrad.codegen.opt import Opt
@@ -43,7 +43,6 @@ def assign_to_contiguous(assign:UOp, target:UOp, src:UOp):
   if target is not t and target.op_in_backward_slice_with_self(Ops.SHRINK):
     # base already realized: copy src only if it reads from the same buffer (overlapping read/write hazard)
     if t.op is Ops.CONTIGUOUS: return assign.replace(src=(target, src.contiguous())) if t in src.toposort() else None
-    if t.op is Ops.CONST: raise RuntimeError("setitem target must be a writable view backed by a buffer")
     mops: list[UOp] = []
     while target.op in GroupOp.Movement:
       mops.append(target)
@@ -313,7 +312,7 @@ DEVICE_MAX_BUFS = {"METAL": 31, "WEBGPU": 8} # TODO: get from device?
 def limit_bufs(ctx:IndexingContext, root:UOp):
   if (device:=root._device) is None: return None # no device, index related calculations
   device = device if isinstance(device, str) else device[0].split(":")[0]
-  if not (MAX_BUFS:=getenv("MAX_KERNEL_BUFFERS", DEVICE_MAX_BUFS.get(device, 0))): return None
+  if not (MAX_BUFS:=MAX_KERNEL_BUFFERS.value or DEVICE_MAX_BUFS.get(device, 0)): return None
 
   bufs: set[UOp] = set()
   def gate_input(u:UOp):
@@ -325,7 +324,7 @@ def limit_bufs(ctx:IndexingContext, root:UOp):
   if len(bufs) > MAX_BUFS - 1: # NOTE: this -1 is for the output buffer
     srcs = []
     for s in root.src:
-      if s.op in GroupOp.Elementwise:
+      if s.op in GroupOp.Elementwise and s._device is not None:
         # Insert bufferize: all AxisType.REDUCE before bufferize are AxisType.LOOP
         orig_ranges, end_ranges = s.ranges, [x.replace(arg=(next(ctx.range_idx), AxisType.LOOP)) if x.op is Ops.RANGE else x for x in s.ranges]
         s = s.substitute(dict(zip(orig_ranges, end_ranges))).bufferize(*end_ranges, arg=BufferizeOpts(device=s.device)).index(*orig_ranges)
@@ -555,7 +554,7 @@ def tag_uop(ctx:tuple[list[UOp], set[UOp]], x:UOp):
   return x.replace(tag=(len(ctx[0])-1,))
 add_tags = pm_gate_kernel_sink+PatternMatcher([
   # don't tag BUFFERs, they are global
-  (UPat(GroupOp.All-{Ops.PARAM, Ops.CONST, Ops.DEVICE, Ops.UNIQUE, Ops.LUNIQUE, Ops.DEFINE_VAR, Ops.BIND, Ops.CALL, Ops.END,
+  (UPat(GroupOp.All-{Ops.PARAM, Ops.CONST, Ops.DEVICE, Ops.UNIQUE, Ops.LUNIQUE, Ops.DEFINE_VAR, Ops.BIND, Ops.END,
                      Ops.MSTACK, Ops.MSELECT, Ops.RANGE}.union(GroupOp.Movement), name="x"), tag_uop),
   (UPat({Ops.MSTACK, Ops.MSELECT}, name="x"), lambda ctx,x: None if all(s.op is Ops.PARAM for s in x.src) else tag_uop(ctx, x)),
 ])
@@ -602,15 +601,15 @@ def get_rangeify_map(sink:UOp) -> dict[UOp, UOp]:
                         name="bufferize to store")
   tsink = graph_rewrite(tsink, pm_gate_kernel_sink+split_kernels, ctx=uop_list, bottom_up=True, name="split kernels")
 
-  # if a kernel depends on a buffer, and that buffer is later assigned to, make the assign depend on the kernel's assign
-  kernel_assign: dict[UOp, UOp] = {}
+  # WAR deps: if kernel U reads buffer S, and S is also written by another kernel, S's write must wait for U to finish
+  afters = [u for u in tsink.toposort() if u.op is Ops.AFTER]
+  kernel_assign: dict[UOp, UOp] = {u.buf_uop:u for u in afters}
   assign_rep: dict[UOp, UOp] = {}
-  for u in tsink.toposort():
-    if u.op is not Ops.AFTER: continue
-    kernel_assign[u.buf_uop] = u
+  for u in afters:
     for s in u.src[1].src:
       # TODO: this is probably broken for MSELECT/MSTACK
       if s.op not in {Ops.BUFFER, Ops.PARAM} or s is u.buf_uop or (a:=kernel_assign.get(s)) is None: continue
+      if a.src[1] is u.src[1]: continue  # same kernel (multi-output custom kernels)
       if any(x.op is Ops.AFTER and x.buf_uop is s for x in u.toposort()):
         raise RuntimeError(f"cycle detected in graph, kernel for {u.buf_uop} must either depend on AFTER or BUFFER")
       assign_rep[a] = kernel_assign[s] = a.replace(src=a.src+(u,))

tinygrad/tensor.py

@@ -614,14 +614,15 @@ class Tensor(OpMixin):
     print(t.numpy())
     ```
     """
-    if not dtypes.is_float(dtype := to_dtype(dtype or dtypes.default_float)): raise ValueError(f"rand only supports float dtypes, got {dtype}")
+    dt = to_dtype(dtype or dtypes.default_float)
+    if not dtypes.is_float(dt): raise ValueError(f"rand only supports float dtypes, got {dt}")
     if not all_int(shape:=argfix(*shape)) or not all(s >= 0 for s in shape): raise ValueError(f"invalid input {shape=}")
     if device is not None and not isinstance(device, str): raise ValueError(f"rand only supports single device, got {device=}")
     device = cast(str, canonicalize_device(device))
 
     # if shape has 0, return zero tensor
-    if (numel := prod(shape)) == 0: return Tensor.zeros(shape, device=device, dtype=dtype, **kwargs)
-    num = ceildiv(numel * dtype.itemsize, 4)
+    if (numel := prod(shape)) == 0: return Tensor.zeros(shape, device=device, dtype=dt, **kwargs)
+    num = ceildiv(numel * dt.itemsize, 4)
 
     # generate per device seeds and rng counter if we haven't seen this device yet
     if device not in Tensor._device_seeds:
@@ -639,14 +640,14 @@ class Tensor(OpMixin):
     bits = Tensor._threefry_random_bits(Tensor._device_seeds[device], counts0, counts1)[:num]
 
     # bitcast to uint with same number of bits
-    _, nmant = dtypes.finfo(dtype)
-    uint_dtype = {1: dtypes.uint8, 2: dtypes.uint16, 4: dtypes.uint32, 8: dtypes.uint64}[dtype.itemsize]
+    _, nmant = dtypes.finfo(dt)
+    uint_dtype = {1: dtypes.uint8, 2: dtypes.uint16, 4: dtypes.uint32, 8: dtypes.uint64}[dt.itemsize]
     bits = bits.bitcast(uint_dtype)
     # only randomize the mantissa bits and set the exponent to 1
-    one = Tensor.ones_like(bits, device=bits.device, dtype=dtype).bitcast(uint_dtype)
-    bits = bits.rshift(dtype.bitsize - nmant).bitwise_or(one)
+    one = Tensor.ones_like(bits, device=bits.device, dtype=dt).bitcast(uint_dtype)
+    bits = bits.rshift(dt.bitsize - nmant).bitwise_or(one)
     # bitcast back to the original dtype and reshape
-    out = bits.bitcast(dtype)[:numel].sub(1).reshape(shape).requires_grad_(kwargs.get("requires_grad"))
+    out = bits.bitcast(dt)[:numel].sub(1).reshape(shape).requires_grad_(kwargs.get("requires_grad"))
     return out.contiguous() if contiguous else out
 
   # ***** creation helper functions *****
@@ -770,8 +771,9 @@ class Tensor(OpMixin):
     print(Tensor.eye(2, 4).numpy())
     ```
     """
-    if n < 0 or ((m := n if m is None else m) < 0): raise ValueError(f"cannot have negative {n=}, {m=}")
-    t = (Tensor.arange(n, device=device).unsqueeze(-1) == Tensor.arange(m, device=device))
+    m_ = n if m is None else m
+    if n < 0 or m_ < 0: raise ValueError(f"cannot have negative {n=}, {m_=}")
+    t = (Tensor.arange(n, device=device).unsqueeze(-1) == Tensor.arange(m_, device=device))
     return t.cast(dtype or dtypes.default_float).requires_grad_(requires_grad)
 
   def _multi_like(self, fxn, *args, **kwargs) -> Tensor:
@@ -1214,6 +1216,26 @@ class Tensor(OpMixin):
     x_dims = [p for p in indices_parsed if not isinstance(p['index'], sint)]
     x = x.reshape(tuple(p['size'] for p in x_dims))
 
+    # basic setitem: construct result with view region replaced by v using arange masks
+    if v is not None and not any(isinstance(p['index'], Tensor) for p in indices_parsed):
+      # broadcast v to getitem shape, reshape to self.ndim (squeeze None dims, unsqueeze int dims — all are size 1)
+      vb = v.cast(self.dtype)._broadcast_to(x.shape)
+      vb = vb.reshape(tuple(1 if isinstance(p['index'], sint) else p['size'] for p in indices_parsed if p['index'] is not None))
+      # undo movement ops per-dim and build boolean mask
+      per_dim = []
+      for d, m in enumerate(mops):
+        (s, e), st = m['boundary'], abs(m['stride'])
+        if st != 1 and vb.shape[d] > 1:  # un-stride: interleave with zeros
+          vb = vb.unsqueeze(d+1)
+          vb = vb.pad_to(tuple(st if j == d+1 else None for j in range(vb.ndim)))
+          vb = vb.reshape(vb.shape[:d] + (vb.shape[d]*vb.shape[d+1],) + vb.shape[d+2:])
+          vb = vb.shrink_to(tuple(e-s if j == d else None for j in range(self.ndim)))
+        idx = Tensor.arange(self.shape[d], device=self.device).reshape([1]*d + [self.shape[d]] + [1]*(self.ndim - d - 1))
+        per_dim.append((idx >= s) & (idx < e) & (((e-1-idx) if m['stride'] < 0 else (idx-s)) % st == 0))
+      vb = vb.flip(tuple(d for d, m in enumerate(mops) if m['stride'] < 0))
+      vb = vb.pad(tuple((m['boundary'][0], self.shape[d] - m['boundary'][1]) for d, m in enumerate(mops)))
+      return (functools.reduce(lambda a, b: a & b, per_dim) if per_dim else Tensor(True, dtype=dtypes.bool, device=self.device)).where(vb, self)
+
     # tensor indexing
     if tops := [(d, p) for d, p in enumerate(x_dims) if isinstance(p['index'], Tensor)]:
       dims, tensors, masks = [d for d, _ in tops], cast(list[Tensor], [p['index'] for _, p in tops]), []
@@ -1309,10 +1331,13 @@ class Tensor(OpMixin):
       if is_disk: raise RuntimeError("advanced setitem is not supported for DISK tensors")
       if not isinstance(v, Tensor): v = Tensor(v, device=self.device, dtype=self.dtype)
       self.assign(self._getitem(indices, v))
-    else: # basic setitem
-      if is_disk: self[indices].assign(v)
-      else:
-        self[indices].assign(v).realize()
+    elif is_disk or self.uop.is_realized: # basic setitem, self is realized. TODO: disk uop.base is a COPY and not realized
+      self[indices].assign(v)
+    else: # basic setitem, self is not realized
+      if not isinstance(v, Tensor): v = Tensor(v, device=self.device, dtype=self.dtype)
+      # __iadd__/__isub__ on unrealized views creates a no-op ASSIGN; unwrap to get the computed value
+      if v.uop.op is Ops.ASSIGN: v = v._apply_uop(lambda x: x.src[1])
+      self.replace(self._getitem(indices, v))
 
   def __delitem__(self, indices) -> None:
     raise TypeError("Tensor does not support deleting items")
@@ -3879,10 +3904,7 @@ class Tensor(OpMixin):
     print(t.dtype, t.numpy())
     ```
     """
-    if (dt:=to_dtype(dtype)) in {dtypes.uint8, dtypes.uint16} and dtypes.is_float(self.dtype):
-      # NOTE: values within the int32 range and outside the unsigned dtype range will cause values to wrap around
-      return self._apply_uop(UOp.cast, dtype=dtypes.int32)._apply_uop(UOp.cast, dtype=dt)
-    return self if self.dtype == dt else self._apply_uop(UOp.cast, dtype=dt)
+    return self if self.dtype == (dt:=to_dtype(dtype)) else self._apply_uop(UOp.cast, dtype=dt)
 
   def bitcast(self, dtype:DTypeLike) -> Tensor:
     """

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()
 
-  # machine instruction: arg=Inst object, tag=register assignment
+  # INS is a machine instruction
   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, Ops.INS}):
+    if self.op in GroupOp.ALU.union({Ops.CAST, Ops.COPY, Ops.NOOP, Ops.GROUP, Ops.SINK, Ops.ALLREDUCE, Ops.STORE}):
       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),
 
-  # machine instruction (ISel output)
+  # assembly instruction
   (UPat(Ops.INS), lambda: True),
 
   # INDEX (2-arg and 3-arg with bool gate)