something

move image index
Merge branch 'master' into image_no_vec
2026-08-15 14:58:28 +00:00 · 2026-05-08 17:32:30 -07:00 · 2026-05-08 17:27:34 -07:00 · 2026-05-08 16:32:08 -07:00 · 2026-05-08 16:31:10 -07:00 · 2026-05-08 16:27:55 -07:00
450 changed files with 140645 additions and 125285 deletions
@@ -225,14 +225,12 @@ runs:
      if: inputs.amd == 'true' && runner.os == 'Linux'
      shell: bash
      run: |
-        cargo build --release --manifest-path ./extra/remu/Cargo.toml
-        sudo ln -sf ${{ github.workspace }}/extra/remu/target/release/libremu.so /usr/local/lib/libremu.so
        sudo tee --append /etc/ld.so.conf.d/rocm.conf <<'EOF'
          /opt/rocm/lib
          /opt/rocm/lib64
        EOF
        sudo ldconfig
-    - name: Setup AMD comgr+remu (macOS)
+    - name: Setup AMD comgr (macOS)
      if: inputs.amd == 'true' && runner.os == 'macOS'
      shell: bash
      run: |
@@ -240,7 +238,6 @@ runs:
        curl -s -H "Authorization: token $GH_TOKEN" curl -s https://api.github.com/repos/tinygrad/amdcomgr_dylib/releases/latest | \
          jq -r '.assets[] | select(.name == "libamd_comgr.dylib").browser_download_url' | \
          sudo xargs curl -fL -o /usr/local/lib/libamd_comgr.dylib
-        cargo build --release --manifest-path ./extra/remu/Cargo.toml

    # **** gpuocelot ****

@@ -33,12 +33,8 @@ jobs:
      uses: ./.github/actions/setup-tinygrad
      with:
        key: 'autogen'
-        opencl: 'true'
        amd: 'true'
-        cuda: 'true'
        llvm: 'true'
-        webgpu: 'true'
-        mesa: 'true'
        pydeps: 'pyyaml mako'
    - name: Install autogen support packages
      run: sudo apt-get install -y --no-install-recommends libclang-20-dev llvm-20-dev hip-dev libusb-1.0-0-dev libdrm-dev
@@ -48,7 +44,7 @@ jobs:
        python3 -c "from tinygrad.runtime.autogen import opencl"
        python3 -c "from tinygrad.runtime.autogen import cuda, nvrtc, nvjitlink, nv_570, nv_580, nv"
        python3 -c "from tinygrad.runtime.autogen import comgr_3, hsa, hip, amd_gpu, sqtt, rocprof, amdgpu_kd, amdgpu_drm"
-        python3 -c "from tinygrad.runtime.autogen.am import am, pm4_soc15, pm4_nv, sdma_4_0_0, sdma_5_0_0, sdma_6_0_0, smu_v13_0_0, smu_v13_0_6, smu_v13_0_12, smu_v14_0_2"
+        python3 -c "from tinygrad.runtime.autogen.am import *"
        python3 -c "from tinygrad.runtime.autogen import libc, kfd, io_uring, ib, pci, vfio"
        python3 -c "from tinygrad.runtime.autogen import llvm"
        python3 -c "from tinygrad.runtime.autogen import webgpu"
@@ -57,6 +53,8 @@ jobs:
        python3 -c "from tinygrad.runtime.autogen import mesa"
        python3 -c "from tinygrad.runtime.autogen import avcodec"
        python3 -c "from tinygrad.runtime.autogen import llvm_qcom"
+        python3 -c "from tinygrad.runtime.autogen import mlx5"
+        python3 -c "from tinygrad.runtime.autogen import ggml_common"
        REGEN=1 python3 -c "from tinygrad.runtime.autogen import libclang"
    - name: Check for differences
      run: |
@@ -51,6 +51,36 @@ jobs:
    - name: openpilot compile3 0.10.1 driving_vision
      run: FLOAT16=1 DEV=CL IMAGE=1 python3.11 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/720392c9a5b986981fdbed1bb8c47a6c5573a50e/selfdrive/modeld/models/driving_vision.onnx

+  # TODO: reenable when not flaky
+  #testframeworkpytest:
+  #  name: framework pytest
+  #  env:
+  #    CI: ""
+  #    CAPTURE_PROCESS_REPLAY: "0"
+  #  runs-on: [self-hosted, framework]
+  #  timeout-minutes: 10
+  #  defaults:
+  #    run:
+  #      shell: bash -e -o pipefail {0}
+  #  if: github.repository_owner == 'tinygrad'
+  #  steps:
+  #  - name: Checkout Code
+  #    uses: actions/checkout@v6
+  #  - name: setup python environment
+  #    run: |
+  #      rm -rf /tmp/tinygrad_pytest_ci
+  #      uv venv /tmp/tinygrad_pytest_ci
+  #      source /tmp/tinygrad_pytest_ci/bin/activate
+  #      uv pip install .[testing]
+  #  - name: setup staging db
+  #    run: |
+  #      echo "CACHEDB=/tmp/pytest-db-ci.db" >> $GITHUB_ENV
+  #      rm -f /tmp/pytest-db-ci*
+  #  - name: Run pytest -nauto
+  #    run: |
+  #      source /tmp/tinygrad_pytest_ci/bin/activate
+  #      pytest -nauto --durations=20
+
  testmacbenchmark:
    name: Mac Benchmark
    env:
@@ -101,10 +131,10 @@ jobs:
      run: DEV=METAL python3.11 test/opt/test_tensor_cores.py
    - name: Test AMX tensor cores
      run: |
-        DEBUG=2 DEV=CPU CPU_LLVM=0 AMX=1 python3.11 test/opt/test_tensor_cores.py
-        DEBUG=2 DEV=CPU CPU_LLVM=1 AMX=1 python3.11 test/opt/test_tensor_cores.py
-        DEBUG=2 DEV=CPU CPU_LLVM=0 AMX=1 python3.11 test/opt/test_gen_float4.py TestFloat4.test_float4_multidim_amx TestFloat4.test_float4_multidim_unaligned_load_amx
-        DEBUG=2 DEV=CPU CPU_LLVM=1 AMX=1 python3.11 test/opt/test_gen_float4.py TestFloat4.test_float4_multidim_amx TestFloat4.test_float4_multidim_unaligned_load_amx
+        DEBUG=2 DEV=CPU AMX=1 python3.11 test/opt/test_tensor_cores.py
+        DEBUG=2 DEV=CPU:LLVM AMX=1 python3.11 test/opt/test_tensor_cores.py
+        DEBUG=2 DEV=CPU AMX=1 python3.11 test/opt/test_gen_float4.py TestFloat4.test_float4_multidim_amx TestFloat4.test_float4_multidim_unaligned_load_amx
+        DEBUG=2 DEV=CPU:LLVM AMX=1 python3.11 test/opt/test_gen_float4.py TestFloat4.test_float4_multidim_amx TestFloat4.test_float4_multidim_unaligned_load_amx
    - name: Run Tensor Core GEMM (float)
      run: DEBUG=2 SHOULD_USE_TC=1 python3.11 extra/gemm/simple_matmul.py
    - name: Run Tensor Core GEMM (half)
@@ -160,7 +190,7 @@ jobs:
        path: |
          onnx_inference_speed.csv
    - name: Run process replay tests
-      run: cp test/external/process_replay/process_replay.py ./process_replay.py && git fetch origin master && git -c advice.detachedHead=false checkout origin/master && PYTHONPATH=. python3.11 process_replay.py
+      uses: ./.github/actions/process-replay

  testusbgpu:
    name: UsbGPU Benchmark
@@ -185,17 +215,19 @@ jobs:
        PYTHONPATH=. ./extra/hcq/hcq_smi.py amd kill_pids
        PYTHONPATH=. ./extra/hcq/hcq_smi.py nv kill_pids
    - name: UsbGPU boot time
-      run: sudo -E PYTHONPATH=. GMMU=0 DEBUG=2 AM_RESET=1 DEV=AMD AMD_IFACE=USB time python3.11 test/test_tiny.py TestTiny.test_plus
+      run: sudo -E PYTHONPATH=. GMMU=0 DEBUG=2 AM_RESET=1 DEV=USB+AMD time python3.11 test/test_tiny.py TestTiny.test_plus
    - name: UsbGPU tiny tests
-      run: sudo -E PYTHONPATH=. GMMU=0 DEV=AMD AMD_IFACE=USB python3.11 test/test_tiny.py
+      run: sudo -E PYTHONPATH=. GMMU=0 DEV=USB+AMD python3.11 test/test_tiny.py
    - name: UsbGPU copy speeds
-      run: sudo -E PYTHONPATH=. GMMU=0 DEV=AMD AMD_IFACE=USB python3.11 test/external/external_test_usb_asm24.py TestDevCopySpeeds
+      run: sudo -E PYTHONPATH=. GMMU=0 DEV=USB+AMD python3.11 test/external/external_test_usb_asm24.py TestDevCopySpeeds
    #- name: UsbGPU openpilot test
-    #  run: sudo -E PYTHONPATH=. GMMU=0 DEV=AMD AMD_IFACE=USB GRAPH_ONE_KERNEL=1 python3.11 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/9118973ed03c1ae1d40cf69a29507ec2cc78efd7/selfdrive/modeld/models/supercombo.onnx
+    #  run: sudo -E PYTHONPATH=. GMMU=0 DEV=USB+AMD GRAPH_ONE_KERNEL=1 python3.11 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/9118973ed03c1ae1d40cf69a29507ec2cc78efd7/selfdrive/modeld/models/supercombo.onnx
+    - name: UsbGPU (USB4/TB) install script
+      run: PYTHONPATH=. sh extra/setup_tinygpu_osx.sh
    - name: UsbGPU (USB4/TB) boot time
-      run: PYTHONPATH=. DEBUG=3 DEV=NV NV_IFACE=PCI NV_NAK=1 time python3.11 test/test_tiny.py TestTiny.test_plus
+      run: PYTHONPATH=. DEBUG=3 DEV=PCI+NV:NAK time python3.11 test/test_tiny.py TestTiny.test_plus
    - name: UsbGPU (USB4/TB) tiny tests
-      run: PYTHONPATH=. DEV=NV NV_IFACE=PCI NV_NAK=1 python3.11 test/test_tiny.py
+      run: PYTHONPATH=. DEV=PCI+NV:NAK python3.11 test/test_tiny.py

  testnvidiabenchmark:
    name: tinybox green Benchmark
@@ -237,7 +269,7 @@ jobs:
    - name: Test tensor cores
      run: |
        DEV=NV ALLOW_TF32=1 python3 test/opt/test_tensor_cores.py
-        DEV=NV NV_PTX=1 ALLOW_TF32=1 python3 test/opt/test_tensor_cores.py
+        DEV=NV:PTX ALLOW_TF32=1 python3 test/opt/test_tensor_cores.py
    - name: Run Tensor Core GEMM (CUDA)
      run: |
        DEV=CUDA SHOULD_USE_TC=1 HALF=1 DEBUG=2 python3 extra/gemm/simple_matmul.py
@@ -245,7 +277,7 @@ jobs:
        DEV=CUDA SHOULD_USE_TC=1 ALLOW_TF32=1 DEBUG=2 ATOL=2e-2 python3 extra/gemm/simple_matmul.py
        DEV=CUDA SHOULD_USE_TC=1 FP8E4M3=1 DEBUG=2 python3 extra/gemm/simple_matmul.py
    - name: Run Tensor Core GEMM (PTX)
-      run: DEV=NV NV_PTX=1 SHOULD_USE_TC=1 HALF=1 DEBUG=2 python3 extra/gemm/simple_matmul.py
+      run: DEV=NV:PTX SHOULD_USE_TC=1 HALF=1 DEBUG=2 python3 extra/gemm/simple_matmul.py
    - name: Run Tensor Core GEMM (NV)
      run: DEV=NV SHOULD_USE_TC=1 HALF=1 DEBUG=2 python3 extra/gemm/simple_matmul.py
    - name: Test DEV=NV
@@ -293,7 +325,7 @@ jobs:
        path: |
          onnx_inference_speed.csv
    - name: Run process replay tests
-      run: cp test/external/process_replay/process_replay.py ./process_replay.py && git fetch origin master && git -c advice.detachedHead=false checkout origin/master && PYTHONPATH=. python3 process_replay.py
+      uses: ./.github/actions/process-replay

  testmorenvidiabenchmark:
    name: tinybox green Training Benchmark
@@ -328,7 +360,7 @@ jobs:
    #   run: DEV=NV M_START=12 M_STOP=20 M_STEP=1 N_START=6 N_STOP=10 N_STEP=1 K_START=28 K_STOP=36 K_STEP=1 HALF=1 TC_OPT=2 python3 ./extra/gemm/fuzz_matmul.py
    # TODO: too slow
    # - name: Fuzz Padded Tensor Core GEMM (PTX)
-    #   run: DEV=NV NV_PTX=1 M_START=12 M_STOP=20 M_STEP=1 N_START=6 N_STOP=10 N_STEP=1 K_START=28 K_STOP=36 K_STEP=1 HALF=1 TC_OPT=2 python3 ./extra/gemm/fuzz_matmul.py
+    #   run: DEV=NV:PTX M_START=12 M_STOP=20 M_STEP=1 N_START=6 N_STOP=10 N_STEP=1 K_START=28 K_STOP=36 K_STEP=1 HALF=1 TC_OPT=2 python3 ./extra/gemm/fuzz_matmul.py
    - name: HEVC Decode Benchmark
      run: VALIDATE=1 MAX_FRAMES=100 ASSERT_FPS=1400 JITBEAM=1 DEV=NV PYTHONPATH=. python3 extra/hevc/decode.py
    - name: Train MNIST
@@ -355,7 +387,7 @@ jobs:
      # TODO: remove BERT_LAYERS once scheduler is fast
      run: BENCHMARK_LOG=bert_10steps_6gpu DEV=NV CAPTURE_PROCESS_REPLAY=0 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=72 GPUS=6 BERT_LAYERS=2 MODEL=bert python3 examples/mlperf/model_train.py
    - name: Run process replay tests
-      run: cp test/external/process_replay/process_replay.py ./process_replay.py && git fetch origin master && git -c advice.detachedHead=false checkout origin/master && PYTHONPATH=. python3 process_replay.py
+      uses: ./.github/actions/process-replay

  testamdbenchmark:
    name: tinybox red Benchmark
@@ -410,11 +442,11 @@ jobs:
    #    LD_PRELOAD="/opt/rocm/lib/libhsa-runtime64.so" HSA=1 BIG=2 TORCHCUDA=1 python3 test/speed/external_test_speed_v_torch.py
    - name: Test speed vs theoretical
      run: DEV=AMD IGNORE_BEAM_CACHE=1 CCACHE=0 BEAM_DEBUG=1 DEBUG=1 python -m pytest -rA test/external/speed_v_theoretical.py --durations=20
-    - name: Test tensor cores AMD_LLVM=0
-      run: DEV=AMD AMD_LLVM=0 python3 test/opt/test_tensor_cores.py
+    - name: Test tensor cores (no LLVM)
+      run: DEV=AMD python3 test/opt/test_tensor_cores.py
    # TODO: this is flaky
-    # - name: Test tensor cores AMD_LLVM=1
-    #   run: DEV=AMD AMD_LLVM=1 python3 test/opt/test_tensor_cores.py
+    # - name: Test tensor cores AMD:LLVM
+    #   run: DEV=AMD:LLVM python3 test/opt/test_tensor_cores.py
    - name: Run Tensor Core GEMM (AMD)
      run: |
        DEV=AMD SHOULD_USE_TC=1 BFLOAT16=1 DEBUG=2 python3 extra/gemm/simple_matmul.py
@@ -467,7 +499,7 @@ jobs:
    - name: Run GPT2 w HALF/BEAM
      run: BENCHMARK_LOG=gpt2_half_beam DEV=AMD HALF=1 JITBEAM=2 IGNORE_BEAM_CACHE=1 python3 examples/gpt2.py --count 10 --temperature 0 --timing
    - name: Run process replay tests
-      run: cp test/external/process_replay/process_replay.py ./process_replay.py && git fetch origin master && git -c advice.detachedHead=false checkout origin/master && PYTHONPATH=. python3 process_replay.py
+      uses: ./.github/actions/process-replay

  testmoreamdbenchmark:
    name: tinybox red Training Benchmark
@@ -503,6 +535,8 @@ jobs:
        rm -f /tmp/staging.db /tmp/staging.db-shm /tmp/staging.db-wal
    - name: reset process replay
      run: test/external/process_replay/reset.py
+    - name: Test GPU crash recovery
+      run: DEV=AMD python3 -m pytest -rA test/external/external_test_gpu_crash.py
    - name: Train MNIST
      run: time PYTHONPATH=. DEV=AMD TARGET_EVAL_ACC_PCT=96.0 python3 examples/beautiful_mnist.py
    - name: Run 10 CIFAR training steps
@@ -522,7 +556,7 @@ jobs:
    #- name: Test full tinyfs load
    #  run: TINYFS_ENDPOINT=10.0.52.11:6767 PYTHONPATH=. python extra/tinyfs/fetch_file.py --hash d734f5e3be9f1e9d863bfaa4fc6c1ef2 --len 175866113 --dest mapping.json --check
    - name: Run process replay tests
-      run: cp test/external/process_replay/process_replay.py ./process_replay.py && git fetch origin master && git -c advice.detachedHead=false checkout origin/master && PYTHONPATH=. python3 process_replay.py
+      uses: ./.github/actions/process-replay

  testmlperfamdbenchmark:
    name: tinybox red MLPerf Benchmark
@@ -568,7 +602,7 @@ jobs:
      # TODO: remove BERT_LAYERS once scheduler is fast
      run: BENCHMARK_LOG=bert_10steps_6gpu DEV=AMD CAPTURE_PROCESS_REPLAY=0 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=72 GPUS=6 BERT_LAYERS=2 MODEL=bert python3 examples/mlperf/model_train.py
    - name: Run process replay tests
-      run: cp test/external/process_replay/process_replay.py ./process_replay.py && git fetch origin master && git -c advice.detachedHead=false checkout origin/master && PYTHONPATH=. python3 process_replay.py
+      uses: ./.github/actions/process-replay

  testqualcommbenchmark:
    name: comma Benchmark
@@ -590,8 +624,10 @@ jobs:
      run: test/external/process_replay/reset.py
    - name: openpilot compile3 0.11.0 driving_vision
      run: BENCHMARK_LOG=openpilot_0_11_0_vision PYTHONPATH="." ASSERT_MIN_STEP_TIME=17 DEV=QCOM FLOAT16=1 IMAGE=1 NOLOCALS=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.11.0/selfdrive/modeld/models/driving_vision.onnx
+    - name: openpilot compile3 0.11.0 driving_vision (from pickle)
+      run: BENCHMARK_LOG=openpilot_0_11_0_vision_run_pickle RUN_PICKLE=1 PYTHONPATH="." ASSERT_MIN_STEP_TIME=17 DEV=QCOM taskset -c 4-7 python3 examples/openpilot/compile3.py
    - name: IR3 openpilot compile3 0.11.0 driving_vision
-      run: BENCHMARK_LOG=ir3_openpilot_0_11_0_vision PYTHONPATH="." ASSERT_MIN_STEP_TIME=17 DEV=QCOM QCOM_IR3=1 FLOAT16=1 IMAGE=1 NOLOCALS=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.11.0/selfdrive/modeld/models/driving_vision.onnx
+      run: BENCHMARK_LOG=ir3_openpilot_0_11_0_vision PYTHONPATH="." ASSERT_MIN_STEP_TIME=17 DEV=QCOM:IR3 FLOAT16=1 IMAGE=1 NOLOCALS=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.11.0/selfdrive/modeld/models/driving_vision.onnx
    - name: openpilot compile3 0.11.0 driving_policy
      run: BENCHMARK_LOG=openpilot_0_11_0_policy PYTHONPATH="." ASSERT_MIN_STEP_TIME=3 DEV=QCOM FLOAT16=1 IMAGE=1 NOLOCALS=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.11.0/selfdrive/modeld/models/driving_policy.onnx
    - name: openpilot compile3 0.11.0 dmonitoring
@@ -609,11 +645,11 @@ jobs:
        # generate quantized weights
        ln -s /data/home/tiny/tinygrad/extra/datasets/imagenet extra/datasets/imagenet
        ln -s /data/home/tiny/tinygrad/testsig-*.so .
-        PYTHONPATH=. CC=clang-19 DEV=CPU CPU_LLVM=0 QUANT=1 CNT=0 python3 examples/test_onnx_imagenet.py https://github.com/xamcat/mobcat-samples/raw/refs/heads/master/onnx_runtime/InferencingSample/InferencingSample/mobilenetv2-7.onnx /tmp/model.quant.onnx
+        PYTHONPATH=. CC=clang-19 DEV=CPU QUANT=1 CNT=0 python3 examples/test_onnx_imagenet.py https://github.com/xamcat/mobcat-samples/raw/refs/heads/master/onnx_runtime/InferencingSample/InferencingSample/mobilenetv2-7.onnx /tmp/model.quant.onnx
        # benchmark on DSP with NOOPT=1, the devectorizer has issues
        PYTHONPATH=. CC=clang-19 DEV=DSP NOOPT=1 CNT=2 DEBUG=2 python3 examples/test_onnx_imagenet.py /tmp/model.quant.onnx
    - name: Run process replay tests
-      run: cp test/external/process_replay/process_replay.py ./process_replay.py && git fetch origin master && git -c advice.detachedHead=false checkout origin/master && PYTHONPATH=. python3 process_replay.py
+      uses: ./.github/actions/process-replay

  testcommausbgpubenchmark:
    name: UsbGPU Benchmark (comma)
@@ -632,9 +668,11 @@ jobs:
        echo "CACHEDB=/tmp/staging.db" >> $GITHUB_ENV
        rm -f /tmp/staging.db /tmp/staging.db-shm /tmp/staging.db-wal
    - name: openpilot compile3 0.10.1 driving_vision
-      run: BENCHMARK_LOG=usbgpu_openpilot_0_10_1_vision PYTHONPATH="." GMMU=0 DEV=AMD AMD_LLVM=1 AMD_IFACE=USB ASSERT_MIN_STEP_TIME=50 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/720392c9a5b986981fdbed1bb8c47a6c5573a50e/selfdrive/modeld/models/driving_vision.onnx
+      run: BENCHMARK_LOG=usbgpu_openpilot_0_10_1_vision PYTHONPATH="." GMMU=0 DEV=USB+AMD:LLVM ASSERT_MIN_STEP_TIME=50 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/720392c9a5b986981fdbed1bb8c47a6c5573a50e/selfdrive/modeld/models/driving_vision.onnx
    - name: openpilot load_pickle 0.10.1 driving_vision
-      run: BENCHMARK_LOG=usbgpu_openpilot_0_10_1_vision_load_pickle PYTHONPATH="." GMMU=0 DEV=AMD AMD_IFACE=USB ASSERT_MIN_LOAD_TIME=15 python3 examples/openpilot/load_pickle.py
+      run: BENCHMARK_LOG=usbgpu_openpilot_0_10_1_vision_load_pickle PYTHONPATH="." GMMU=0 DEV=USB+AMD ASSERT_MIN_LOAD_TIME=15 python3 examples/openpilot/load_pickle.py
+    - name: openpilot run_pickle 0.10.1 driving_vision
+      run: BENCHMARK_LOG=usbgpu_openpilot_0_10_1_vision_run_pickle RUN_PICKLE=1 PYTHONPATH="." GMMU=0 DEV=USB+AMD ASSERT_MIN_STEP_TIME=50 python3 examples/openpilot/compile3.py

  testreddriverbenchmark:
    name: AM Benchmark
@@ -674,11 +712,13 @@ jobs:
      run: time DEBUG=3 DEV=AMD AM_RESET=1 python3 test/test_tiny.py TestTiny.test_plus
    - name: Test driver warm start time
      run: time DEBUG=3 DEV=AMD python3 test/test_tiny.py TestTiny.test_plus
+    - name: Test GPU crash recovery
+      run: DEV=AMD python3 -m pytest -rA test/external/external_test_gpu_crash.py
    # Fails on 9070
    # - name: Test tensor cores
    #   run: |
-    #     DEV=AMD AMD_LLVM=0 python3 test/test_linearizer.py test/opt/test_tensor_cores.py
-    #     DEV=AMD AMD_LLVM=1 python3 test/test_linearizer.py test/opt/test_tensor_cores.py
+    #     DEV=AMD python3 test/test_linearizer.py test/opt/test_tensor_cores.py
+    #     DEV=AMD:LLVM python3 test/test_linearizer.py test/opt/test_tensor_cores.py
    #     DEV=AMD SHOULD_USE_TC=1 BFLOAT16=1 DEBUG=2 python3 extra/gemm/simple_matmul.py
    - name: Run Tensor Core GEMM (AMD)
      run: DEV=AMD SHOULD_USE_TC=1 HALF=1 DEBUG=2 ATOL=2e-2 python3 extra/gemm/simple_matmul.py
@@ -702,11 +742,11 @@ jobs:
        pkill -f 'extra/remote/serve.py' || true
        PYTHONPATH=. python3 extra/remote/serve.py 6482 &
        sleep 1
-        DEBUG=2 PYTHONPATH=. REMOTE=127.0.0.1:6482 AM_RESET=1 DEV=AMD AMD_IFACE=PCI python3 test/test_tiny.py
-        DEBUG=2 PYTHONPATH=. REMOTE=127.0.0.1:6482 AM_RESET=1 DEV=AMD AMD_AQL=1 AMD_IFACE=PCI python3 test/test_tiny.py
+        DEBUG=2 PYTHONPATH=. REMOTE=127.0.0.1:6482 AM_RESET=1 DEV=PCI+AMD python3 test/test_tiny.py
+        DEBUG=2 PYTHONPATH=. REMOTE=127.0.0.1:6482 AM_RESET=1 DEV=PCI+AMD AMD_AQL=1 python3 test/test_tiny.py
        pkill -f 'extra/remote/serve.py' || true
    - name: Run process replay tests
-      run: cp test/external/process_replay/process_replay.py ./process_replay.py && git fetch origin master && git -c advice.detachedHead=false checkout origin/master && PYTHONPATH=. python3 process_replay.py
+      uses: ./.github/actions/process-replay

  testgreendriverbenchmark:
    name: NV Benchmark
@@ -769,4 +809,4 @@ jobs:
        DEBUG=2 PYTHONPATH=. REMOTE=127.0.0.1:6483 DEV=NV python3 test/test_tiny.py
        pkill -f 'extra/remote/serve.py' || true
    - name: Run process replay tests
-      run: cp test/external/process_replay/process_replay.py ./process_replay.py && git fetch origin master && git -c advice.detachedHead=false checkout origin/master && PYTHONPATH=. python3 process_replay.py
+      uses: ./.github/actions/process-replay
@@ -66,9 +66,7 @@ jobs:
          PR="$GITHUB_WORKSPACE/pr"
          pip install tabulate $BASE
          cp "$BASE/sz.py" .
-          echo "loc_content<<EOF" >> "$GITHUB_ENV"
-          python sz.py "$BASE" "$PR" >> "$GITHUB_ENV"
-          echo "EOF" >> "$GITHUB_ENV"
+          python sz.py "$BASE" "$PR" > loc_content.txt
      - name: Comment Code Line Diff
        continue-on-error: false
        uses: marocchino/sticky-pull-request-comment@v3
@@ -77,7 +75,7 @@ jobs:
          ignore_empty: true
          skip_unchanged: true
          recreate: true
-          message: ${{ env.loc_content }}
+          path: loc_content.txt

  rebase:
    name: Core Library Line Difference
@@ -1,7 +1,7 @@
 name: Unit Tests
 env:
  # increment this when downloads substantially change to avoid the internet
-  CACHE_VERSION: '18'
+  CACHE_VERSION: '19'
  CAPTURE_PROCESS_REPLAY: 1
  GH_TOKEN: ${{ secrets.GITHUB_TOKEN }}
  PYTHONPATH: ${{ github.workspace }}
@@ -29,9 +29,9 @@ jobs:
        deps: testing_unit
        llvm: 'true'
    - name: Speed Test
-      run: DEV=CPU CPU_LLVM=1 THREADS=0 python3 test/speed/external_test_speed_v_torch.py
+      run: DEV=CPU:LLVM THREADS=0 python3 test/speed/external_test_speed_v_torch.py
    - name: Speed Test (BEAM=2)
-      run: BEAM=2 DEV=CPU CPU_LLVM=1 THREADS=0 python3 test/speed/external_test_speed_v_torch.py
+      run: BEAM=2 DEV=CPU:LLVM THREADS=0 python3 test/speed/external_test_speed_v_torch.py

  docs:
    name: Docs
@@ -83,7 +83,7 @@ jobs:
      run: DEBUG=100 python3 -c "from tinygrad import Tensor; N = 1024; a, b = Tensor.rand(N, N), Tensor.rand(N, N); c = (a.reshape(N, 1, N) * b.T.reshape(1, N, N)).sum(axis=2); print((c.numpy() - (a.numpy() @ b.numpy())).mean())"
    - name: Compile EfficientNet to C and test it
      run: |
-        DEV=CPU CPU_LLVM=0 python examples/compile_efficientnet.py > recognize.c
+        DEV=CPU python examples/compile_efficientnet.py > recognize.c
        clang -O2 recognize.c -lm -o recognize
        cat test/models/efficientnet/Chicken.jpg | ./recognize | grep cock

@@ -114,11 +114,11 @@ jobs:
    - name: Test one op in torch tests
      run: DEBUG=2 python3 extra/torch_backend/torch_tests.py TestTinyBackendPRIVATEUSE1.test_unary_log_tiny_float32
    - name: Test Ops with TINY_BACKEND
-      run: DEV=CPU CPU_LLVM=1 LLVMOPT=0 TINY_BACKEND=1 python3 -m pytest -n auto test/backend/test_ops.py --durations=20
+      run: DEV=CPU:LLVM LLVMOPT=0 TINY_BACKEND=1 python3 -m pytest -n auto test/backend/test_ops.py --durations=20
    - name: Test in-place operations on views
      run: TORCH_DEBUG=1 python3 extra/torch_backend/test_inplace.py
    - name: Test multi-gpu
-      run: DEV=CPU CPU_LLVM=1 GPUS=4 TORCH_DEBUG=1 python3 extra/torch_backend/test_multigpu.py
+      run: DEV=CPU:LLVM GPUS=4 TORCH_DEBUG=1 python3 extra/torch_backend/test_multigpu.py
    - name: Test kernel fusion
      run: python3 extra/torch_backend/test_kernel_fusion.py

@@ -173,45 +173,45 @@ jobs:
        IMAGE=1 DEV=PYTHON python3 test/backend/test_ops.py TestOps.test_simple_conv2d
    - name: Test emulated METAL tensor cores
      run: |
-        DEBUG=2 EMULATE=METAL FORWARD_ONLY=1 DEV=PYTHON python3 test/backend/test_ops.py TestOps.test_big_gemm
-        DEBUG=2 EMULATE=METAL FORWARD_ONLY=1 DEV=PYTHON python3 test/opt/test_tensor_cores.py
+        DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::METAL python3 test/backend/test_ops.py TestOps.test_big_gemm
+        DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::METAL python3 test/opt/test_tensor_cores.py
    - name: Test emulated AMX tensor cores
-      run: DEBUG=2 AMX=1 EMULATE=AMX FORWARD_ONLY=1 DEV=PYTHON python3 test/backend/test_ops.py TestOps.test_gemm
+      run: DEBUG=2 AMX=1 FORWARD_ONLY=1 DEV=PYTHON::AMX python3 test/backend/test_ops.py TestOps.test_gemm
    - name: Test emulated AMD tensor cores
      run: |
-        DEBUG=2 EMULATE=AMD FORWARD_ONLY=1 DEV=PYTHON N=16 HALF=1 ACC_HALF=0 python3 ./extra/gemm/simple_matmul.py
-        DEBUG=2 EMULATE=AMD FORWARD_ONLY=1 DEV=PYTHON N=64 HALF=1 ACC_HALF=0 python3 ./extra/gemm/simple_matmul.py
-        DEBUG=2 EMULATE=AMD FORWARD_ONLY=1 DEV=PYTHON N=16 HALF=1 ACC_HALF=1 ATOL=1e-3 python3 ./extra/gemm/simple_matmul.py
-        DEBUG=2 EMULATE=AMD FORWARD_ONLY=1 DEV=PYTHON N=64 HALF=1 ACC_HALF=1 ATOL=1e-3 python3 ./extra/gemm/simple_matmul.py
-        DEBUG=2 EMULATE=AMD FORWARD_ONLY=1 DEV=PYTHON python3 test/opt/test_tensor_cores.py
+        DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::gfx1100 N=16 HALF=1 ACC_HALF=0 python3 ./extra/gemm/simple_matmul.py
+        DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::gfx1100 N=64 HALF=1 ACC_HALF=0 python3 ./extra/gemm/simple_matmul.py
+        DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::gfx1100 N=16 HALF=1 ACC_HALF=1 ATOL=1e-3 python3 ./extra/gemm/simple_matmul.py
+        DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::gfx1100 N=64 HALF=1 ACC_HALF=1 ATOL=1e-3 python3 ./extra/gemm/simple_matmul.py
+        DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::gfx1100 python3 test/opt/test_tensor_cores.py
    - name: Test emulated AMD MFMA tensor cores
      run: |
-        DEBUG=2 EMULATE=AMD_MFMA FORWARD_ONLY=1 DEV=PYTHON N=64 HALF=1 ACC_HALF=0 python3 ./extra/gemm/simple_matmul.py
-        DEBUG=2 EMULATE=AMD_MFMA FORWARD_ONLY=1 DEV=PYTHON python3 test/opt/test_tensor_cores.py
+        DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::gfx950 N=64 HALF=1 ACC_HALF=0 python3 ./extra/gemm/simple_matmul.py
+        DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::gfx950 python3 test/opt/test_tensor_cores.py
    - name: Test emulated AMD RDNA4 tensor cores
      run: |
-        DEBUG=2 EMULATE=AMD_RDNA4 FORWARD_ONLY=1 DEV=PYTHON N=16 HALF=1 ACC_HALF=0 python3 ./extra/gemm/simple_matmul.py
-        DEBUG=2 EMULATE=AMD_RDNA4 FORWARD_ONLY=1 DEV=PYTHON N=64 HALF=1 ACC_HALF=0 python3 ./extra/gemm/simple_matmul.py
-        DEBUG=2 EMULATE=AMD_RDNA4 FORWARD_ONLY=1 DEV=PYTHON N=16 HALF=1 ACC_HALF=1 ATOL=1e-3 python3 ./extra/gemm/simple_matmul.py
-        DEBUG=2 EMULATE=AMD_RDNA4 FORWARD_ONLY=1 DEV=PYTHON N=64 HALF=1 ACC_HALF=1 ATOL=1e-3 python3 ./extra/gemm/simple_matmul.py
-        DEBUG=2 EMULATE=AMD_RDNA4 FORWARD_ONLY=1 DEV=PYTHON python3 test/opt/test_tensor_cores.py
+        DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::gfx1201 N=16 HALF=1 ACC_HALF=0 python3 ./extra/gemm/simple_matmul.py
+        DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::gfx1201 N=64 HALF=1 ACC_HALF=0 python3 ./extra/gemm/simple_matmul.py
+        DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::gfx1201 N=16 HALF=1 ACC_HALF=1 ATOL=1e-3 python3 ./extra/gemm/simple_matmul.py
+        DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::gfx1201 N=64 HALF=1 ACC_HALF=1 ATOL=1e-3 python3 ./extra/gemm/simple_matmul.py
+        DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::gfx1201 python3 test/opt/test_tensor_cores.py
    - name: Test emulated CUDA tensor cores
      run: |
-        DEBUG=2 EMULATE=CUDA FORWARD_ONLY=1 DEV=PYTHON python3 test/backend/test_ops.py TestOps.test_gemm_fp16
-        DEBUG=2 EMULATE=CUDA ALLOW_TF32=1 FORWARD_ONLY=1 DEV=PYTHON python3 test/backend/test_ops.py TestOps.test_gemm
-        DEBUG=2 EMULATE=CUDA_SM75 FORWARD_ONLY=1 DEV=PYTHON python3 test/backend/test_ops.py TestOps.test_gemm_fp16
-        DEBUG=2 EMULATE=CUDA_SM89 ALLOW_TF32=1 FORWARD_ONLY=1 DEV=PYTHON python3 test/opt/test_tensor_cores.py
+        DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::sm_80 python3 test/backend/test_ops.py TestOps.test_gemm_fp16
+        DEBUG=2 ALLOW_TF32=1 FORWARD_ONLY=1 DEV=PYTHON::sm_80 python3 test/backend/test_ops.py TestOps.test_gemm
+        DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::sm_75 python3 test/backend/test_ops.py TestOps.test_gemm_fp16
+        DEBUG=2 ALLOW_TF32=1 FORWARD_ONLY=1 DEV=PYTHON::sm_89 python3 test/opt/test_tensor_cores.py
    - name: Test emulated INTEL OpenCL tensor cores
-      run: DEBUG=2 EMULATE=INTEL FORWARD_ONLY=1 DEV=PYTHON HALF=1 N=64 python3 ./extra/gemm/simple_matmul.py
+      run: DEBUG=2 FORWARD_ONLY=1 DEV=PYTHON::INTEL HALF=1 N=64 python3 ./extra/gemm/simple_matmul.py
    - name: Test emulated AMX tensor cores
-      run: DEBUG=2 AMX=1 EMULATE=AMX FORWARD_ONLY=1 DEV=PYTHON python3 test/opt/test_tensor_cores.py
+      run: DEBUG=2 AMX=1 FORWARD_ONLY=1 DEV=PYTHON::AMX python3 test/opt/test_tensor_cores.py
    - name: Test device flop counts
      run: |
-        DEBUG=2 EMULATE=METAL DEV=PYTHON python3 ./test/null/test_uops_stats.py TestUOpsStatsMatmulHalf
-        DEBUG=2 EMULATE=AMD DEV=PYTHON python3 ./test/null/test_uops_stats.py TestUOpsStatsMatmulHalf
-        DEBUG=2 EMULATE=CUDA DEV=PYTHON python3 ./test/null/test_uops_stats.py TestUOpsStatsMatmulHalf
-        DEBUG=2 EMULATE=INTEL DEV=PYTHON python3 ./test/null/test_uops_stats.py TestUOpsStatsMatmulHalf
-        DEBUG=2 AMX=1 EMULATE=AMX DEV=PYTHON python3 ./test/null/test_uops_stats.py TestUOpsStats.test_simple_matmul
+        DEBUG=2 DEV=PYTHON::METAL python3 ./test/null/test_uops_stats.py TestUOpsStatsMatmulHalf
+        DEBUG=2 DEV=PYTHON::gfx1100 python3 ./test/null/test_uops_stats.py TestUOpsStatsMatmulHalf
+        DEBUG=2 DEV=PYTHON::sm_80 python3 ./test/null/test_uops_stats.py TestUOpsStatsMatmulHalf
+        DEBUG=2 DEV=PYTHON::INTEL python3 ./test/null/test_uops_stats.py TestUOpsStatsMatmulHalf
+        DEBUG=2 AMX=1 DEV=PYTHON::AMX python3 ./test/null/test_uops_stats.py TestUOpsStats.test_simple_matmul

  linter:
    name: Linters
@@ -270,7 +270,7 @@ jobs:
    - name: Run Clip tests for SD MLPerf on NULL backend
      run: DEV=NULL python -m pytest -n=auto test/external/mlperf_stable_diffusion/external_test_models.py::TestOpenClip --durations=20
    - name: Run AMD emulated BERT training on NULL backend
-      run: EMULATE=AMD_RDNA4 DEV=NULL NULL_ALLOW_COPYOUT=1 CAPTURE_PROCESS_REPLAY=0 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=66 GPUS=1 BERT_LAYERS=2 MODEL=bert python3 examples/mlperf/model_train.py
+      run: DEV=NULL::gfx1201 NULL_ALLOW_COPYOUT=1 CAPTURE_PROCESS_REPLAY=0 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=66 GPUS=1 BERT_LAYERS=2 MODEL=bert python3 examples/mlperf/model_train.py
    # TODO: support fake weights
    #- name: Run LLaMA 7B on 4 fake devices
    #  run: DEV=NULL python3 examples/llama.py --gen 1 --size 7B --shard 4 --prompt "Hello." --count 3 --temperature 0 --timing
@@ -333,7 +333,7 @@ jobs:
        deps: testing_unit
        python-version: '3.14'
    - name: Test SPEC=2
-      run: SPEC=2 pytest --maxfail=10 -n auto --durations=30 test/unit test/backend test/opt --ignore test/backend/test_custom_kernel.py --ignore test/unit/test_hashing.py --timeout 60 -k "not test_setitem_big" --splits 2 --group ${{ matrix.group }}
+      run: SPEC=2 pytest --maxfail=10 -n auto --durations=30 test/unit test/backend test/opt --ignore test/backend/test_custom_kernel.py --ignore test/unit/test_hashing.py --timeout 60 -k "not test_setitem_big" -k "not test_conv2d_ceildiv_edge_case" --splits 2 --group ${{ matrix.group }}

  fuzzing:
    name: Fuzzing
@@ -417,13 +417,13 @@ jobs:
          llvm: 'true'
      - name: Test openpilot model kernel count and gate usage
        run: |
-          ALLOWED_KERNEL_COUNT=123 ALLOWED_READ_IMAGE=1486 ALLOWED_GATED_READ_IMAGE=17 FLOAT16=1 DEV=CL IMAGE=1 python examples/openpilot/compile3.py https://gitlab.com/commaai/openpilot-lfs.git/gitlab-lfs/objects/cf6376aa9a090f0da26c280ef69eabf9bbdd51d1faac9ed392919c3db69be916
+          ALLOWED_KERNEL_COUNT=123 ALLOWED_READ_IMAGE=1486 ALLOWED_GATED_READ_IMAGE=18 FLOAT16=1 DEV=CL IMAGE=1 python examples/openpilot/compile3.py https://gitlab.com/commaai/openpilot-lfs.git/gitlab-lfs/objects/cf6376aa9a090f0da26c280ef69eabf9bbdd51d1faac9ed392919c3db69be916
      - name: Test openpilot CL compile fp16
        run: FLOAT16=1 DEV=CL IMAGE=1 python examples/openpilot/compile3.py https://gitlab.com/commaai/openpilot-lfs.git/gitlab-lfs/objects/cf6376aa9a090f0da26c280ef69eabf9bbdd51d1faac9ed392919c3db69be916
      - name: Test openpilot CL compile fp32 (test correctness)
        run: DEV=CL IMAGE=1 SELFTEST=1 python examples/openpilot/compile3.py https://github.com/haraschax/filedump/raw/refs/heads/master/driving_vision_fp32.onnx
      - name: Test openpilot LLVM compile fp16
-        run: IMAGE=1 FLOAT16=1 DEV=CPU CPU_LLVM=1 python examples/openpilot/compile3.py https://gitlab.com/commaai/openpilot-lfs.git/gitlab-lfs/objects/cf6376aa9a090f0da26c280ef69eabf9bbdd51d1faac9ed392919c3db69be916
+        run: IMAGE=1 FLOAT16=1 DEV=CPU:LLVM python examples/openpilot/compile3.py https://gitlab.com/commaai/openpilot-lfs.git/gitlab-lfs/objects/cf6376aa9a090f0da26c280ef69eabf9bbdd51d1faac9ed392919c3db69be916
      - name: Run process replay tests
        uses: ./.github/actions/process-replay

@@ -445,15 +445,15 @@ jobs:
          python-version: '3.12'
          llvm: 'true'
      - name: Test ONNX (CPU)
-        run: DEV=CPU CPU_LLVM=0 python -m pytest -n=auto test/external/external_test_onnx_backend.py --durations=20
+        run: DEV=CPU python -m pytest -n=auto test/external/external_test_onnx_backend.py --durations=20
      - name: Test ONNX (LLVM)
-        run: DEV=CPU CPU_LLVM=1 python -m pytest -n=auto test/external/external_test_onnx_backend.py --durations=20
+        run: DEV=CPU:LLVM python -m pytest -n=auto test/external/external_test_onnx_backend.py --durations=20
      - name: Test ONNX Runner (CPU)
-        run: DEV=CPU CPU_LLVM=0 python3 test/external/external_test_onnx_runner.py
+        run: DEV=CPU python3 test/external/external_test_onnx_runner.py
      - name: Test Additional ONNX Ops (CPU)
-        run: DEV=CPU CPU_LLVM=0 python3 test/external/external_test_onnx_ops.py
+        run: DEV=CPU python3 test/external/external_test_onnx_ops.py
      - name: Test Quantize ONNX
-        run: DEV=CPU CPU_LLVM=0 python3 test/backend/test_quantize_onnx.py
+        run: DEV=CPU python3 test/backend/test_quantize_onnx.py
      - name: Run process replay tests
        uses: ./.github/actions/process-replay

@@ -505,12 +505,14 @@ jobs:
        with:
          key: apps_llm
      - name: Test 1B LLM (llama)
-        run: echo "What's a male chicken called? Answer with only one word." | MAX_BUFFER_SIZE=0 python3 -m tinygrad.apps.llm --model llama3.2:1b | tee /dev/stderr | grep -i rooster
+        run: echo "What's a male chicken called? Answer with only one word." | MAX_BUFFER_SIZE=0 python3 -m tinygrad.llm --model llama3.2:1b | tee /dev/stderr | grep -i rooster
      - name: Test 1B LLM (llama q4)
-        run: echo "What's a male chicken called? Answer with only one word." | MAX_BUFFER_SIZE=0 python3 -m tinygrad.apps.llm --model llama3.2:1b-q4 | tee /dev/stderr | grep -i rooster
+        run: echo "What's a male chicken called? Answer with only one word." | MAX_BUFFER_SIZE=0 python3 -m tinygrad.llm --model llama3.2:1b-q4 | tee /dev/stderr | grep -i rooster
+      - name: Test 1B LLM (qwen3.5)
+        run: echo "What's a male chicken called? Answer with only one word." | MAX_BUFFER_SIZE=0 python3 -m tinygrad.llm --model qwen3.5:0.8b | tee /dev/stderr | grep -i rooster
      - name: Test 1B LLM (qwen)
        # NOTE: qwen is dumb and only knows about female chickens
-        run: echo "What's a female chicken called? Answer with only one word." | MAX_BUFFER_SIZE=0 python3 -m tinygrad.apps.llm --model qwen3:0.6b | tee /dev/stderr | grep -i hen
+        run: echo "What's a female chicken called? Answer with only one word." | MAX_BUFFER_SIZE=0 python3 -m tinygrad.llm --model qwen3:0.6b | tee /dev/stderr | grep -i hen

 # ****** Models Tests ******

@@ -529,11 +531,11 @@ jobs:
          opencl: 'true'
          llvm: 'true'
      - name: Test models (llvm)
-        run: DEV=CPU CPU_LLVM=1 python -m pytest -n=auto test/models --durations=20
+        run: DEV=CPU:LLVM python -m pytest -n=auto test/models --durations=20
      - name: Test models (opencl)
        run: DEV=CL python -m pytest -n=auto test/models --durations=20
      - name: Test models (cpu)
-        run: DEV=CPU CPU_LLVM=0 python -m pytest -n=auto test/models --durations=20
+        run: DEV=CPU python -m pytest -n=auto test/models --durations=20
      - name: Run process replay tests
        uses: ./.github/actions/process-replay

@@ -572,11 +574,11 @@ jobs:
        pydeps: "pillow"
        llvm: "true"
    - name: Test LLVM=1 DEVECTORIZE=0
-      run: DEV=CPU CPU_LLVM=1 DEVECTORIZE=0 python3 -m pytest -n auto test/test_tiny.py test/backend/test_ops.py
+      run: DEV=CPU:LLVM DEVECTORIZE=0 python3 -m pytest -n auto test/test_tiny.py test/backend/test_ops.py
    - name: Test LLVM=1 DEVECTORIZE=0 for model
-      run: DEV=CPU CPU_LLVM=1 DEVECTORIZE=0 python3 test/models/test_efficientnet.py
+      run: DEV=CPU:LLVM DEVECTORIZE=0 python3 test/models/test_efficientnet.py
    - name: Test DEV=CPU DEVECTORIZE=0
-      run: DEV=CPU CPU_LLVM=0 DEVECTORIZE=0 python3 -m pytest -n auto test/test_tiny.py test/backend/test_ops.py
+      run: DEV=CPU DEVECTORIZE=0 python3 -m pytest -n auto test/test_tiny.py test/backend/test_ops.py

  testdsp:
    name: Linux (DSP)
@@ -641,9 +643,7 @@ jobs:
    runs-on: ubuntu-24.04
    timeout-minutes: 20
    env:
-      DEV: AMD
-      PYTHON_REMU: 1
-      MOCKGPU: 1
+      DEV: MOCKKFD+AMD
    steps:
      - name: Checkout Code
        uses: actions/checkout@v6
@@ -667,30 +667,30 @@ jobs:
      - name: Install rocprof-trace-decoder
        run: sudo PYTHONPATH="." ./extra/sqtt/install_rocprof_decoder.py
      - name: Run AMD renderer tests
-        run: AMD_LLVM=0 python -m pytest -n=auto test/amd/ --durations 20
-      - name: Run AMD renderer tests (AMD_LLVM=1)
-        run: AMD_LLVM=1 python -m pytest -n=auto test/amd/ --durations 20
+        run: python -m pytest -n=auto test/amd/ --durations 20
+      - name: Run AMD renderer tests (AMD:LLVM)
+        run: DEV=MOCKKFD+AMD:LLVM python -m pytest -n=auto test/amd/ --durations 20
      - name: Run SQTT profiling tests
        run: PROFILE=1 SQTT=1 python3 -m pytest -n=auto test/amd/test_sqtt_profiler.py
      - name: Run AMD emulated tests on NULL backend
        env:
          AMD: 0
        run: |
-          PYTHONPATH=. DEV=NULL EMULATE=AMD python extra/mmapeak/mmapeak.py
-          PYTHONPATH=. DEV=NULL EMULATE=AMD_CDNA4 python3 -m pytest -n=auto test/testextra/test_tk.py test/backend/test_asm_gemm.py
-      - name: Run ASM matmul on MOCKGPU
-        run: PYTHONPATH="." DEV=AMD MOCKGPU=1 N=256 python3 extra/gemm/amd_asm_matmul.py
+          PYTHONPATH=. DEV=NULL:HIP:gfx1100 python extra/mmapeak/mmapeak.py
+          PYTHONPATH=. DEV=NULL:HIP:gfx950 python3 -m pytest -n=auto test/testextra/test_tk.py test/backend/test_asm_gemm.py
+      - name: Run matmul on MOCKKFD
+        run: |
+          PYTHONPATH="." DEV=MOCKKFD+AMD N=256 python3 extra/gemm/amd_asm_matmul.py
+          PYTHONPATH="." DEV=MOCKKFD+AMD N=256 python3 extra/gemm/amd_copy_matmul.py
      - name: Run LLVM test
-        run: AMD_LLVM=1 python test/device/test_amd_llvm.py
+        run: DEV=MOCKKFD+AMD:LLVM python test/device/test_amd_llvm.py

  testmockam:
    name: Linux (am)
    runs-on: ubuntu-24.04
    timeout-minutes: 15
    env:
-      DEV: AMD
-      MOCKGPU: 1
-      AMD_IFACE: PCI
+      DEV: MOCKPCI+AMD
    steps:
      - name: Checkout Code
        uses: actions/checkout@v6
@@ -702,13 +702,13 @@ jobs:
          amd: 'true'
      - name: Run test_tiny on MOCKAM
        run: python test/test_tiny.py
-      - name: Run test_tiny on MOCKAM USB
-        run: GMMU=0 AMD_IFACE=USB python test/test_tiny.py
-      - name: Run test_hcq on MOCKAM
+      - name: Run test_tiny on MOCKUSB
+        run: GMMU=0 DEV=MOCKUSB+AMD python test/test_tiny.py
+      - name: Run test_hcq on MOCKPCI
        run: python -m pytest test/device/test_hcq.py
-      - name: Run disk copy tests on MOCKAM
+      - name: Run disk copy tests on MOCKPCI
        run: python -m pytest test/unit/test_disk_tensor.py -k test_copy_from_disk
-      - name: Run test_tiny on MOCKAM Remote
+      - name: Run test_tiny on MOCKPCI Remote
        run: |
          python extra/remote/serve.py 6667 &
          sleep 2
@@ -720,17 +720,14 @@ jobs:
      fail-fast: false
      matrix:
        backend: [amd, amdllvm]
-        arch: [rdna3, rdna4, cdna4]
+        arch: [gfx1100, gfx1201, gfx950]

    name: Linux (${{ matrix.backend }} ${{ matrix.arch }})
    runs-on: ubuntu-22.04
    timeout-minutes: 15
    env:
-      DEV: AMD
-      MOCKGPU: 1
-      MOCKGPU_ARCH: ${{ matrix.arch }}
+      DEV: MOCKKFD+AMD:${{ matrix.backend == 'amdllvm' && 'LLVM' || '' }}:${{ matrix.arch }}
      SKIP_SLOW_TEST: 1
-      AMD_LLVM: ${{ matrix.backend == 'amdllvm' && '1' || matrix.backend != 'amdllvm' && '0' }}
    steps:
      - name: Checkout Code
        uses: actions/checkout@v6
@@ -764,7 +761,6 @@ jobs:
    runs-on: ubuntu-22.04
    timeout-minutes: 20
    env:
-      MOCKGPU: 1
      FORWARD_ONLY: 1
    steps:
      - name: Checkout Code
@@ -777,7 +773,7 @@ jobs:
          cuda: 'true'
          ocelot: 'true'
      - name: Set env
-        run: printf "${{ matrix.backend == 'ptx' && 'DEV=CUDA\nCUDA_PTX=1' || matrix.backend == 'nv' && 'DEV=NV\nSKIP_SLOW_TEST=1' }}" >> $GITHUB_ENV
+        run: printf "${{ matrix.backend == 'ptx' && 'DEV=MOCK+CUDA:PTX' || matrix.backend == 'nv' && 'DEV=MOCK+NV\nSKIP_SLOW_TEST=1' }}" >> $GITHUB_ENV
      - name: Check Device.DEFAULT and print some source
        run: |
          python3 -c "from tinygrad import Device; assert Device.DEFAULT in ['CUDA','NV'], Device.DEFAULT"
@@ -811,7 +807,7 @@ jobs:
          llvm: ${{ matrix.backend == 'llvm' || matrix.backend == 'lvp' }}
          mesa: ${{ matrix.backend == 'lvp' && 'true' }}
      - name: Set env
-        run: printf "${{ matrix.backend == 'llvm' && 'DEV=CPU\nCPU_LLVM=1' || matrix.backend == 'cpu' && 'DEV=CPU\nCPU_LLVM=0\nCPU_COUNT=2' || matrix.backend == 'opencl' && 'DEV=CL' || matrix.backend == 'lvp' && 'DEV=CPU\nCPU_LVP=1' }}" >> $GITHUB_ENV
+        run: printf "${{ matrix.backend == 'llvm' && 'DEV=CPU:LLVM' || matrix.backend == 'cpu' && 'DEV=CPU\nCPU_COUNT=2' || matrix.backend == 'opencl' && 'DEV=CL' || matrix.backend == 'lvp' && 'DEV=CPU:LVP' }}" >> $GITHUB_ENV
      - name: Check Device.DEFAULT and print some source
        run: |
          python3 -c "from tinygrad import Device; assert Device.DEFAULT in ['CPU','CL'], Device.DEFAULT"
@@ -862,25 +858,19 @@ jobs:
      run: DEV=METAL TRANSCENDENTAL=2 python -m pytest -n=auto test/backend/test_ops.py::TestOps::test_sin test/backend/test_ops.py::TestOps::test_cos test/backend/test_ops.py::TestOps::test_tan test/backend/test_ops.py::TestOps::test_exp test/backend/test_ops.py::TestOps::test_log --durations=20
    - name: Run pytest (amd)
      env:
-        MOCKGPU: 1
-        DEV: AMD
-        AMD_LLVM: 0
+        DEV: MOCKKFD+AMD
        FORWARD_ONLY: 1
      run: |
        python3 -m pytest -n=auto test/device/test_hcq.py test/test_tiny.py --durations=20
    - name: Run pytest (amd with llvm backend)
      env:
-        MOCKGPU: 1
-        DEV: AMD
-        AMD_LLVM: 1
+        DEV: "MOCKKFD+AMD:LLVM"
        FORWARD_ONLY: 1
      run: |
        python -m pytest -n=auto test/device/test_hcq.py test/test_tiny.py test/device/test_amd_llvm.py --durations=20
    - name: Run pytest (ptx)
      env:
-        MOCKGPU: 1
-        NV_PTX: 1
-        DEV: NV
+        DEV: "MOCK+NV:PTX"
        FORWARD_ONLY: 1
        # TODO: failing due to library loading error
        CAPTURE_PROCESS_REPLAY: 0
@@ -945,7 +935,7 @@ jobs:
          llvm: ${{ matrix.backend == 'llvm' || matrix.backend == 'lvp' }}
          mesa: ${{ matrix.backend == 'lvp' && 'true' }}
      - name: Set env
-        run: printf "${{ matrix.backend == 'llvm' && 'DEV=CPU\nCPU_LLVM=1' || matrix.backend == 'cpu' && 'DEV=CPU\nCPU_LLVM=0\nCPU_COUNT=2' || matrix.backend == 'metal' && 'DEV=METAL' || matrix.backend == 'lvp' && 'DEV=CPU\nCPU_LVP=1' }}" >> $GITHUB_ENV
+        run: printf "${{ matrix.backend == 'llvm' && 'DEV=CPU:LLVM' || matrix.backend == 'cpu' && 'DEV=CPU\nCPU_COUNT=2' || matrix.backend == 'metal' && 'DEV=METAL' || matrix.backend == 'lvp' && 'DEV=CPU:LVP' }}" >> $GITHUB_ENV
      - name: Check Device.DEFAULT and print some source
        run: |
          python -c "from tinygrad import Device; assert Device.DEFAULT == {'LLVM':'CPU','LVP':'CPU'}.get(x:='${{ matrix.backend }}'.upper(), x), Device.DEFAULT"
@@ -955,8 +945,8 @@ jobs:
      - name: Run process replay tests
        uses: ./.github/actions/process-replay
      - name: Run macOS-specific unit test
-        if: matrix.backend == 'cpu'
-        run: python3 -m pytest test/unit/test_disk_tensor.py::TestDiskTensor::test_copy_to_cpu_not_truncated
+        if: matrix.backend == 'llvm'
+        run: python3 -m pytest test/unit/test_disk_tensor.py::TestDiskTensor::test_copy_to_cpu_not_truncated test/unit/test_cpu.py

 # ****** Windows Tests ******

@@ -980,7 +970,7 @@ jobs:
          pydeps: ${{ matrix.backend == 'webgpu' && 'dawn-python' || '' }}
      - name: Set env
        shell: bash
-        run:  printf "${{ matrix.backend == 'llvm' && 'DEV=CPU\nCPU_LLVM=1' || matrix.backend == 'cpu' && 'DEV=CPU\nCPU_LLVM=0\nCPU_COUNT=2' || matrix.backend == 'webgpu' && 'DEV=WEBGPU'}}" >> $GITHUB_ENV
+        run:  printf "${{ matrix.backend == 'llvm' && 'DEV=CPU:LLVM' || matrix.backend == 'cpu' && 'DEV=CPU\nCPU_COUNT=2' || matrix.backend == 'webgpu' && 'DEV=WEBGPU'}}" >> $GITHUB_ENV
      - name: Run unit tests
        if: matrix.backend=='llvm'
        # test_newton_schulz hits RecursionError
@@ -988,7 +978,7 @@ jobs:
      - name: Run NULL backend tests
        if: matrix.backend=='llvm'
        shell: bash
-        run: CPU=0 CPU_LLVM=0 DEV=NULL python -m pytest -n=auto test/null/ --ignore=test/null/test_elf.py --durations=20
+        run: DEV=NULL python -m pytest -n=auto test/null/ --ignore=test/null/test_elf.py --durations=20
      - name: Run pytest (${{ matrix.backend }})
        shell: bash
        run: |
@@ -1017,7 +1007,7 @@ jobs:
          python-version: '3.12'
      - name: Set env
        shell: bash
-        run: printf "DEV=NULL\nNULL_ALLOW_COPYOUT=1\n${{ matrix.backend == 'ir3' && 'NULL_IR3=1' || matrix.backend == 'nak' && 'NULL_NAK=1' }}" >> $GITHUB_ENV
+        run: printf "NULL_ALLOW_COPYOUT=1\n${{ matrix.backend == 'ir3' && 'DEV=NULL:IR3:a630' || matrix.backend == 'nak' && 'DEV=NULL:NAK:sm_120' }}" >> $GITHUB_ENV
      - name: Run test_ops
        shell: bash
        run: |
@@ -1040,7 +1030,7 @@ jobs:
          python-version: '3.12'
      - name: Set env
        shell: bash
-        run: printf "DEV=NULL\nNULL_ALLOW_COPYOUT=1\nNULL_QCOMCL=1" >> $GITHUB_ENV
+        run: printf "DEV=NULL:QCOMCL:a630\nNULL_ALLOW_COPYOUT=1" >> $GITHUB_ENV
      - name: Run test_ops
        shell: bash
        run: |
@@ -68,3 +68,4 @@ mutants
 .mutmut-cache
 dagre/
 graphlib/
+uv.lock
@@ -1,6 +1,4 @@
 # abstractions2 goes from back to front, here we will go from front to back
-from typing import List
-from tinygrad.helpers import tqdm

 # *****
 # 0. Load mnist on the device
@@ -33,21 +31,21 @@ model(X).sparse_categorical_crossentropy(Y).backward()
 optim.schedule_step()   # this will step the optimizer without running realize

 # *****
-# 3. Create a schedule.
+# 3. Create a schedule (linear uop).

 # The weight Tensors have been assigned to, but not yet realized. Everything is still lazy at this point
 # l1.uop and l2.uop define a computation graph

-from tinygrad.engine.schedule import ExecItem
-schedule: List[ExecItem] = Tensor.schedule(l1, l2)
+from tinygrad.engine.realize import run_linear
+linear = Tensor.schedule_linear(l1, l2)

-print(f"The schedule contains {len(schedule)} items.")
-for si in schedule: print(str(si)[:80])
+print(f"The schedule contains {len(linear.src)} items.")
+for call in linear.src: print(str(call)[:80])

 # *****
-# 4. Lower and run the schedule.
+# 4. Lower and run the schedule (linear uop).

-for si in tqdm(schedule): si.run()
+run_linear(linear)

 # *****
 # 5. Print the weight change
@@ -0,0 +1,253 @@
+# tinygrad allows you to write kernels at many different abstractions levels.
+# This is for RDNA3, but if you don't have one you can run with the emulator
+# PYTHONPATH="." DEV=MOCKPCI+AMD
+
+from tinygrad import Tensor, Context, GlobalCounters, UOp, Device
+from tinygrad.helpers import DEV, DEBUG, getenv
+from tinygrad.uop.ops import AxisType, KernelInfo, Ops
+from tinygrad.dtype import AddrSpace, dtypes
+from tinygrad.runtime.autogen.amd.rdna3.ins import *
+
+def eval_harness(name, tensor, fxn, check=None):
+  print(f"***** {name}")
+  GlobalCounters.reset()
+  with Context(DEBUG=max(DEBUG.value, 2)): out = fxn(tensor).item()
+  assert check is None or abs(out - check) < abs(check) * 1e-3, f"out was wrong {out}, expected {check}, off by {out/check}x"
+  print(f"computed in {GlobalCounters.time_sum_s*1000:.2f} ms, {(a.nbytes()/1e9)/GlobalCounters.time_sum_s:.2f} GB/s")
+  return out
+
+SZ = 256*1024 if DEV.interface.startswith("MOCK") else 1024*1024*1024
+
+def example_2_hip(a:Tensor, correct):
+  GLOBALS = 1024
+  THREADS = 256
+  def hip_reduce_sum(out:UOp, buf:UOp) -> UOp:
+    assert SZ % (GLOBALS * THREADS) == 0
+    CHUNK = SZ // (GLOBALS * THREADS)
+    # NOTE: tinygrad doesn't populate HIP hidden kernargs, so blockDim.x/gridDim.x read as 0.
+    # We hardcode block/grid sizes as constexpr to avoid any dependency on those builtins.
+    code = f"""
+    #include <hip/hip_runtime.h>
+    constexpr unsigned int BLOCK = {THREADS};
+    constexpr unsigned int CHUNK = {CHUNK};
+    extern "C" __global__ void hip_reduce_sum_kernel(float* __restrict__ block_sums, const float* __restrict__ x) {{
+      __shared__ float sdata[BLOCK];
+
+      unsigned int tid = threadIdx.x;
+      unsigned int gid = blockIdx.x * BLOCK + tid;
+
+      // Each thread sums CHUNK consecutive elements from its own region
+      float sum = 0.0f;
+      const float* base = x + gid * CHUNK;
+      #pragma unroll 16
+      for (unsigned int k = 0; k < CHUNK; k++) {{
+        sum += base[k];
+      }}
+
+      sdata[tid] = sum;
+      __syncthreads();
+
+      // Block reduction in shared memory
+      for (unsigned int s = BLOCK / 2; s > 0; s >>= 1) {{
+        if (tid < s) {{
+          sdata[tid] += sdata[tid + s];
+        }}
+        __syncthreads();
+      }}
+
+      // One partial sum per block
+      if (tid == 0) {{
+        block_sums[blockIdx.x] = sdata[0];
+      }}
+    }}"""
+
+    # TODO: remove the need for the compiler here, you should just be able to remove Ops.BINARY
+    from tinygrad.runtime.support.compiler_amd import HIPCCCompiler
+    lib = HIPCCCompiler(Device[Device.DEFAULT].renderer.target.arch, []).compile_cached(code)
+    # the sink specifies the GLOBAL and LOCAL sizes, along with the input buffers and name
+    sink = UOp.sink(UOp.special(GLOBALS, 'gidx0'), UOp.special(THREADS, 'lidx0'), out, buf,
+                    arg=KernelInfo(name="hip_reduce_sum_kernel"))
+    return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=Device.DEFAULT),
+                UOp(Ops.LINEAR, src=(*sink.src, sink)), UOp(Ops.SOURCE, arg=code), UOp(Ops.BINARY, arg=lib)))
+  eval_harness("HIP kernel", a, lambda x: Tensor.empty(GLOBALS).custom_kernel(x, fxn=hip_reduce_sum)[0].sum(), check=correct)
+
+def example_3_custom_uop(a:Tensor, correct):
+  # This GPU has 32 CUs, keep them all busy
+  CU_COUNT = 32
+  def custom_sum(out:UOp, buf:UOp) -> UOp:
+    LCLS = 256
+    buf = buf.reshape(CU_COUNT, -1, LCLS)
+
+    glbl = UOp.range(CU_COUNT, 0, AxisType.GLOBAL)
+    lane = UOp.range(LCLS, 1, AxisType.LOCAL)
+
+    # accumulate the globals into a per lane accumulator
+    reduce_loop = UOp.range(buf.shape[1], 2, AxisType.REDUCE)
+    acc = UOp.placeholder((1,), dtypes.float, slot=6, addrspace=AddrSpace.REG)
+    acc = acc.after(acc.store(0))
+    acc = acc.after(acc[0].store(acc.after(reduce_loop)[0] + buf[glbl, reduce_loop, lane]).end(reduce_loop))
+
+    # store all the per lane accumulators to LOCAL
+    local_accs = UOp.placeholder((LCLS,), dtypes.float, slot=0, addrspace=AddrSpace.LOCAL)
+    local_accs = local_accs.after(local_accs[lane].store(acc[0]).barrier())
+
+    # accumulate LOCALs into a single per CU accumulator
+    late_reduce_loop = UOp.range(LCLS, 3, AxisType.REDUCE)
+    acc2 = UOp.placeholder((1,), dtypes.float, slot=7, addrspace=AddrSpace.REG)
+    acc2 = acc2.after(acc2.store(0))
+    acc2 = acc2.after(acc2[0].store(acc2.after(late_reduce_loop)[0] + local_accs[late_reduce_loop]).end(late_reduce_loop))[0]
+
+    # store (NOTE: since the address doesn't depend on the warp, this will be automatically gated)
+    return out[glbl].store(acc2).end(lane, glbl).sink(arg=KernelInfo(opts_to_apply=()))
+
+  eval_harness("custom UOp kernel", a, lambda x: Tensor.empty(CU_COUNT).custom_kernel(x, fxn=custom_sum)[0].sum(), check=correct)
+
+def example_5_custom_assembly(a:Tensor, correct):
+  # Kernel class copied from amd_asm_matmul
+  class Kernel:
+    def __init__(self): self.instructions, self.labels, self.pos = [], {}, 0
+    def label(self, name): self.labels[name] = self.pos
+    def emit(self, inst, target=None):
+      self.instructions.append(inst)
+      inst._target, inst._pos = target, self.pos
+      self.pos += inst.size()
+      return inst
+    def waitcnt(self, lgkm=None, vm=None):
+      # Wait for memory operations. lgkm=N waits until N lgkm ops remain, vm=N waits until N vmem ops remain.
+      vmcnt, lgkmcnt, expcnt = vm if vm is not None else 63, lgkm if lgkm is not None else 63, 7
+      waitcnt = (expcnt & 0x7) | ((lgkmcnt & 0x3f) << 4) | ((vmcnt & 0x3f) << 10)
+      self.emit(s_waitcnt(simm16=waitcnt))
+    def finalize(self, sink:UOp) -> UOp:
+      for inst in self.instructions:
+        if inst._target is None: continue
+        offset_dwords = (self.labels[inst._target] - inst._pos - inst.size()) // 4
+        if not -32768 <= offset_dwords <= 32767: raise ValueError(f"branch to '{inst._target}' offset {offset_dwords} exceeds simm16 range")
+        inst.simm16 = offset_dwords
+      return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=Device.DEFAULT),
+                                   UOp(Ops.LINEAR, src=tuple([UOp(Ops.INS, arg=x) for x in self.instructions]))))
+
+  CU_COUNT = 32
+  LANES = 64
+  def asm_sum(out:UOp, buf:UOp) -> UOp:
+    V_LANE_ID = 0             # lane_id set on startup
+    S_WORKGROUP_X = 2         # workgroup_id_x
+    S_LOOP_CTR = 3
+    k = Kernel()
+    # mul lane id by 16 for offsets (4 for float, 4 for b128)
+    k.emit(v_mul_lo_u32(v[0], v[V_LANE_ID], 16))
+    k.emit(v_add_nc_u32_e32(v[1], 4096, v[0]))
+    k.emit(v_add_nc_u32_e32(v[2], 4096, v[1]))
+    k.emit(v_add_nc_u32_e32(v[3], 4096, v[2]))
+    # load both addresses
+    k.emit(s_load_b128(sdata=s[4:7], sbase=s[0:1], offset=0x0, soffset=NULL))
+    k.waitcnt(lgkm=0)
+    # offset buffer pointer by workgroup_id_x * chunk_size_bytes
+    k.emit(s_mul_i32(s[S_LOOP_CTR], s[S_WORKGROUP_X], buf.numel()*4//CU_COUNT))
+    k.emit(s_add_u32(s[6], s[6], s[S_LOOP_CTR]))
+    k.emit(s_addc_u32(s[7], s[7], 0))
+    # zero the accumulators
+    k.emit(VOPD(VOPDOp.V_DUAL_MOV_B32, VOPDOp.V_DUAL_MOV_B32, vdstx=v[4], vdsty=v[5], srcx0=0, srcy0=0))
+    k.emit(VOPD(VOPDOp.V_DUAL_MOV_B32, VOPDOp.V_DUAL_MOV_B32, vdstx=v[6], vdsty=v[7], srcx0=0, srcy0=0))
+
+    def emit_loads(base_vreg, reg_len):
+      assert reg_len%4 == 0
+      k.emit(s_clause(simm16=(reg_len//4)-1))
+      for i in range(reg_len//4):
+        offset = i*LANES*16
+        assert offset < 16384
+        k.emit(global_load_b128(vdst=v[base_vreg+i*4:base_vreg+i*4+3], addr=v[offset//4096], saddr=s[6:7], offset=offset%4096))
+      k.emit(s_add_u32(s[6], s[6], reg_len * LANES * 4))
+      k.emit(s_addc_u32(s[7], s[7], 0))
+
+    def tree_reduce_to_4567(base_vreg, reg_len):
+      assert reg_len%4 == 0
+      reg_len //= 4
+      while reg_len > 1:
+        half = reg_len // 2
+        for j in range(half):
+          a, b = base_vreg + j*4, base_vreg + (j+half)*4
+          # v[a+0](bank0) += v[b+2](bank2), v[a+1](bank1) += v[b+3](bank3) — src0 and src1 on different banks
+          k.emit(VOPD(VOPDOp.V_DUAL_ADD_F32, VOPDOp.V_DUAL_ADD_F32, vdstx=v[a], vdsty=v[a+1], srcx0=v[a], vsrcx1=v[b+2], srcy0=v[a+1], vsrcy1=v[b+3]))
+          # v[a+2](bank2) += v[b+0](bank0), v[a+3](bank3) += v[b+1](bank1) — src0 and src1 on different banks
+          k.emit(VOPD(VOPDOp.V_DUAL_ADD_F32, VOPDOp.V_DUAL_ADD_F32, vdstx=v[a+2], vdsty=v[a+3], srcx0=v[a+2], vsrcx1=v[b], srcy0=v[a+3], vsrcy1=v[b+1]))
+        reg_len = half
+      k.emit(VOPD(VOPDOp.V_DUAL_ADD_F32, VOPDOp.V_DUAL_ADD_F32, vdstx=v[4], vdsty=v[5], srcx0=v[4], vsrcx1=v[base_vreg], srcy0=v[5], vsrcy1=v[base_vreg+1]))
+      k.emit(VOPD(VOPDOp.V_DUAL_ADD_F32, VOPDOp.V_DUAL_ADD_F32, vdstx=v[6], vdsty=v[7], srcx0=v[6], vsrcx1=v[base_vreg+2], srcy0=v[7], vsrcy1=v[base_vreg+3]))
+
+    BASE_REG = 8
+    LOAD_UNROLL = 64
+    INNER_UNROLL = 2
+
+    assert buf.numel() % (CU_COUNT*LANES*LOAD_UNROLL*INNER_UNROLL) == 0
+    total_batches = buf.numel()//(CU_COUNT*LANES*LOAD_UNROLL*INNER_UNROLL)
+    k.emit(s_mov_b32(s[S_LOOP_CTR], total_batches-1))
+
+    k.label('LOOP')
+    for _ in range(INNER_UNROLL):
+      emit_loads(BASE_REG, reg_len=LOAD_UNROLL)
+      k.waitcnt(vm=0)
+      tree_reduce_to_4567(BASE_REG, reg_len=LOAD_UNROLL)
+    k.emit(s_sub_u32(s[S_LOOP_CTR], s[S_LOOP_CTR], 1))
+    k.emit(s_cbranch_scc0(), target='LOOP')
+
+    # add into v[4]
+    k.emit(v_add_f32_e32(v[4], v[4], v[5]))
+    k.emit(v_add_f32_e32(v[6], v[6], v[7]))
+    k.emit(v_add_f32_e32(v[4], v[4], v[6]))
+
+    # warp shuffle into v[4] on lane 0 using DPP row_shl within each 16-lane row
+    for shift in [1, 2, 4, 8]:
+      k.emit(v_add_f32_e32(v[4], DPP, v[4], vsrc0=v[4], dpp=0x100 | shift, row_mask=0xf, bank_mask=0xf, bc=1))
+    # combine rows: get lane 16's value to lane 0 via permlanex16
+    k.emit(v_permlanex16_b32(v[5], v[4], 0, 0))
+    k.emit(v_add_f32_e32(v[4], v[4], v[5]))
+
+    # atomic store (only on lane 0)
+    k.emit(s_mov_b32(EXEC_LO, 1))
+    k.emit(v_mov_b32_e32(v[0], 0))
+    k.emit(global_atomic_add_f32(addr=v[0], saddr=s[4:5], data=v[4]))
+
+    k.emit(s_sendmsg(simm16=3))  # DEALLOC_VGPRS
+    k.emit(s_endpgm())
+    return k.finalize(UOp.sink(UOp.special(CU_COUNT, 'gidx0'), UOp.special(LANES, 'lidx0'), out, buf, arg=KernelInfo(name="asm_reduce")))
+
+  out = Tensor.zeros(1,).contiguous().realize()
+  eval_harness("RDNA3 assembly kernel", a, lambda x: out.custom_kernel(x, fxn=asm_sum)[0], check=correct)
+
+if __name__ == "__main__":
+  examples = [int(x) for x in getenv("EXAMPLES", "1,2,3,4,5").split(",")]
+
+  correct = None
+  # First define a Tensor and realize it. We will focus on a 1GB sum kernel on RDNA3
+  a = (Tensor.randn(SZ) if getenv("RAND") else Tensor.ones(SZ)).contiguous().realize()
+
+  if 1 in examples:
+    # *****
+    # This is the high level tinygrad way.
+    # Note that this is split into multiple kernels for speed.
+    correct = eval_harness("basic kernel", a, lambda x: x.sum())
+
+  if 2 in examples:
+    # *****
+    # You can import kernels from CUDA/HIP/Metal.
+    # ChatGPT is great at writing these Kernel
+    example_2_hip(a, correct)
+
+  if 3 in examples:
+    # *****
+    # Now we get to the lower abstraction layers of tinygrad.
+    # You can write a kernel in UOps, and it's 2.5x faster than normal.
+    example_3_custom_uop(a, correct)
+
+  if 4 in examples:
+    # *****
+    # You can also BEAM search stock tinygrad for a faster kernel.
+    # This does even better than all the kernels to date in this simple case.
+    with Context(BEAM=2):
+      eval_harness("BEAMed kernel", a, lambda x: x.sum(), check=correct)
+
+  if 5 in examples:
+    # *****
+    # If you really want to go crazy with speed, you can code in assembly.
+    # There's not too much to gain here over BEAM, but it's a few percent faster.
+    example_5_custom_assembly(a, correct)
@@ -17,15 +17,13 @@ The `UOp` graph specifies the compute in terms of low level tinygrad ops. Not al

 ## Scheduling

-The [scheduler](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/engine/schedule.py) converts the graph of UOps into a list of `ExecItem`. One `ExecItem` is one kernel on the GPU, and the scheduler is responsible for breaking the large compute graph into subgraphs that can fit in a kernel. `ast` specifies what compute to run, and `bufs` specifies what buffers to run it on.
-
-::: tinygrad.engine.schedule.ExecItem
+The [scheduler](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/schedule/__init__.py) converts the graph of UOps into a `LINEAR` UOp whose `src` is a list of `CALL` UOps. One `CALL` is one kernel on the GPU, and the scheduler is responsible for breaking the large compute graph into subgraphs that can fit in a kernel. The `CALL`'s `src[0]` (a `SINK` ast) specifies what compute to run, and the remaining `src` are the buffers to run it on.

 ## Lowering

-The code in [realize](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/engine/realize.py) lowers `ExecItem` by populating its `prg` field with
+The code in [realize](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/engine/realize.py) lowers each `CALL` by compiling its ast into a `PROGRAM` and running it.

-::: tinygrad.engine.realize.run_schedule
+::: tinygrad.engine.realize.run_linear

 There's a ton of complexity hidden behind this, see the `codegen/` directory.

@@ -35,13 +33,7 @@ Then we render the UOps into code with a `Renderer`, then we compile the code to

 ## Execution

-Creating `ExecItem`, which has a run method
-
-::: tinygrad.engine.realize.ExecItem
-    options:
-        members: true
-
-Lists of `ExecItem` can be condensed into a single ExecItem with the Graph API (rename to Queue?)
+`run_linear` walks the `LINEAR` UOp, dispatching each `CALL` to a runner (kernel, copy, view, encdec, or graph).

 ## Runtime

@@ -28,7 +28,7 @@ Transforms the ast into an optimized ast. This is where BEAM search and heuristi

 Transform the optimized ast into a linearized and rendered program.

-::: tinygrad.codegen.get_program
+::: tinygrad.codegen.to_program
    options:
        members: false
        show_labels: false
@@ -53,7 +53,7 @@ Transform the linearized list of UOps into a program, represented as a string.

 Abstracted high level interface to the runtimes.

-::: tinygrad.engine.realize.get_program
+::: tinygrad.engine.realize.to_program
    options:
        members: false
        show_labels: false
@@ -31,26 +31,43 @@ These control the behavior of core tinygrad even when used as a library.
 Variable | Possible Value(s) | Description
 ---|---|---
 DEBUG               | [1-7]      | enable debugging output (operations, timings, speed, generated code and more)
-DEV                 | [AMD, NV, ...] | enable a specific backend
+DEV                 | [AMD, NV, ...] | enable a specific backend, see [below](#dev-variable)
 BEAM                | [#]        | number of beams in kernel beam search
 DEFAULT_FLOAT       | [HALF, ...]| specify the default float dtype (FLOAT32, HALF, BFLOAT16, FLOAT64, ...), default to FLOAT32
-IMAGE               | [1-2]      | enable 2d specific optimizations
+IMAGE               | [1]        | enable 2d specific optimizations
 FLOAT16             | [1]        | use float16 for images instead of float32
-HCQ_VISIBLE_DEVICES | [list[int]]| restricts the HCQ devices that are available. The format is a comma-separated list of identifiers (indexing starts with 0).
 JIT                 | [0-2]      | 0=disabled, 1=[jit enabled](quickstart.md#jit) (default), 2=jit enabled, but graphs are disabled
 VIZ                 | [1]        | 0=disabled, 1=[viz enabled](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/viz)
 ALLOW_TF32          | [1]        | enable TensorFloat-32 tensor cores on Ampere or newer GPUs.
 WEBGPU_BACKEND      | [WGPUBackendType_Metal, ...]          | Force select a backend for WebGPU (Metal, DirectX, OpenGL, Vulkan...)
 CUDA_PATH           | str        | Use `CUDA_PATH/include` for CUDA headers for CUDA and NV backends. If not set, TinyGrad will use `/usr/local/cuda/include`, `/usr/include` and `/opt/cuda/include`.

-## Debug breakdown
+### DEV variable
+
+The `DEV` variable deserves special note due to its more nuanced syntax.
+`DEV` is used to specify the target device, target renderer and target architecture for said device, separated by colons.
+Specifying the renderer and architecture is optional, omitting a preference will cause tinygrad to automatically determine a suitable setting.
+The `DEV` variable may also be used to specify the interface through which to access the device (eg. `PCI`, `USB`). Interfaces may be specified preceding the target triple,
+separated by a plus (eg. `DEV=USB+AMD:LLVM`). Similarly as above, the interface may be omitted. Example usage follows:
+
+`DEV` contents | Interpretation
+--- | ---
+AMD           | use the AMD device
+AMD:LLVM      | use the AMD device with the LLVM renderer
+NV:CUDA:sm_70 | use the NV device with the CUDA renderer targetting sm_70
+AMD::gfx950   | use the AMD device targetting gfx950
+USB+AMD       | use the AMD device over the USB interface
+CPU:LLVM      | use the CPU device with the LLVM renderer
+CPU:LLVM:x86_64,znver2,avx2,-avx512f | use the CPU device with the LLVM renderer, with [additional arch flags](runtime.md#cpu-arch)
+
+### Debug breakdown

 Variable | Value | Description
 ---|---|---
 DEBUG               | >= 1       | Enables debugging and lists devices being used
 DEBUG               | >= 2       | Provides performance metrics for operations, including timing, memory usage, bandwidth for each kernel execution
-DEBUG               | >= 3       | Outputs buffers used for each kernel (shape, dtype and strides) and the applied optimizations at a kernel level
+DEBUG               | >= 3       | Outputs the applied optimizations at a kernel level
 DEBUG               | >= 4       | Outputs the generated kernel code
-DEBUG               | >= 5       | Displays the intermediate representation of the computation UOps (AST)
+DEBUG               | >= 5       | Displays the intermediate representation of the computation UOps
 DEBUG               | >= 6       | Displays the intermediate representation of the computation UOps in a linearized manner, detailing the operation sequence
 DEBUG               | >= 7       | Outputs the assembly code generated for the target hardware
@@ -37,4 +37,4 @@
    options:
        show_signature: false
        separate_signature: false
-::: tinygrad.nn.state.gguf_load
+::: tinygrad.llm.gguf.gguf_load
@@ -4,13 +4,13 @@ tinygrad supports various runtimes, enabling your code to scale across a wide ra

 | Runtime | Description | Compiler Options | Requirements |
 |---------|-------------|------------------|--------------|
-| [NV](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_nv.py) | Provides acceleration for NVIDIA GPUs | nvrtc (default)<br>PTX (`NV_PTX=1`) | Ampere/Ada/Blackwell series GPUs.<br>You can select an interface via `NV_IFACE=(NVK\|PCI)`. See [NV interfaces](#nv-interfaces) for details. |
-| [AMD](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_amd.py) | Provides acceleration for AMD GPUs | LLVM (`AMD_LLVM=1`)<br>HIP/COMGR (`AMD_HIP=1`) | RDNA2 or newer GPUs.<br>You can select an interface via `AMD_IFACE=(KFD\|PCI\|USB)`. See [AMD interfaces](#amd-interfaces) for details. |
+| [NV](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_nv.py) | Provides acceleration for NVIDIA GPUs | nvrtc (default)<br>PTX (`DEV=NV:PTX`) | Ampere/Ada/Blackwell series GPUs.<br>You can select an interface via [the `DEV` variable](env_vars.md#dev-variable). See [NV interfaces](#nv-interfaces) for details. |
+| [AMD](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_amd.py) | Provides acceleration for AMD GPUs | LLVM (`DEV=AMD:LLVM`)<br>HIP/COMGR (`DEV=AMD:HIP`) | CDNA3, CDNA4, RDNA3 or RDNA4 GPUs.<br>You can select an interface via [the `DEV` variable](env_vars.md#dev-variable). See [AMD interfaces](#amd-interfaces) for details. |
 | [QCOM](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_qcom.py) | Provides acceleration for QCOM GPUs | - | 6xx series GPUs |
 | [METAL](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_metal.py) | Utilizes Metal for acceleration on Apple devices | - | M1+ Macs; Metal 3.0+ for `bfloat` support |
-| [CUDA](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_cuda.py) | Utilizes CUDA for acceleration on NVIDIA GPUs | nvrtc (default)<br> PTX (`CUDA_PTX=1`) | NVIDIA GPU with CUDA support |
+| [CUDA](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_cuda.py) | Utilizes CUDA for acceleration on NVIDIA GPUs | nvrtc (default)<br> PTX (`DEV=CUDA:PTX`) | NVIDIA GPU with CUDA support |
 | [CL](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_cl.py) | Accelerates computations using OpenCL on GPUs | - | OpenCL 2.0 compatible device |
-| [CPU](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_cpu.py) | Runs on CPU using the clang or llvm compiler | Clang JIT (default)<br>LLVM IR (`CPU_LLVM=1`) | `clang` compiler in system `PATH` |
+| [CPU](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_cpu.py) | Runs on CPU using the clang or llvm compiler | Clang JIT (default)<br>LLVM IR (`DEV=CPU:LLVM`) | `clang` compiler in system `PATH`<br>You can specify additional arch parameters via [the `DEV` variable](env_vars.md#dev-variable). See [CPU arch](#cpu-arch) for details. |
 | [WEBGPU](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_webgpu.py) | Runs on GPU using the Dawn WebGPU engine (used in Google Chrome) | - | Dawn library installed and discoverable. Binaries: [pydawn v0.3.0](https://github.com/wpmed92/pydawn/releases/tag/v0.3.0) |


@@ -72,10 +72,20 @@ AMD backend supports several interfaces for communicating with devices:
 * `PCI`: uses the [AM driver](developer/am.md)
 * `USB`: USB3 interface for asm24xx chips.

-You can force an interface by setting `AMD_IFACE` to one of these values. In the case of `AMD_IFACE=PCI`, this may unbind your GPU from the amdgpu driver.
+You can force an interface by setting the interface component of [the `DEV` environment variable](env_vars.md#dev-variable) to one of these values. When set to `PCI`, this may unbind your GPU from the amdgpu driver.

 ## NV Interfaces
 NV backend supports several interfaces for communicating with devices:

 * `NVK`: uses the nvidia driver
 * `PCI`: uses the [NV driver](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/support/nv/nvdev.py)
+
+## CPU Arch
+The CPU renderers may be additionally configured using the arch component of [the `DEV` environment variable](env_vars.md#dev-variable).
+CPU arch should be specified as a comma-separated list of parameters, and must contain at least two values: the architecture family (ie. x86_64, arm64, or riscv64) and the cpu type (as accepted by `clang`'s `-march`).
+If native is specified as the cpu type, tinygrad (or delegate compiler) will query the host cpu type. Additional comma-separated values may be specified as follows:
+
+* `AMX`: emit Apple silicon AMX instructions
+
+All other additional values are interpreted as cpu feature flags. When a value is preceded by a `-` character, the corresponding feature flag will be disabled, otherwise the flag will be enabled.
+Note that enabled feature flags should not be preceded by a `+`.
@@ -66,8 +66,8 @@ Elementwise ops operate on a per element basis. They don't change the shape of t
 ::: tinygrad.Tensor.sub
 ::: tinygrad.Tensor.mul
 ::: tinygrad.Tensor.div
-::: tinygrad.Tensor.idiv
 ::: tinygrad.Tensor.mod
+::: tinygrad.Tensor.fmod
 ::: tinygrad.Tensor.bitwise_xor
 ::: tinygrad.Tensor.bitwise_and
 ::: tinygrad.Tensor.bitwise_or
@@ -19,8 +19,8 @@

 ## tinygrad ops

-::: tinygrad.Tensor.schedule_with_vars
-::: tinygrad.Tensor.schedule
+::: tinygrad.Tensor.linear_with_vars
+::: tinygrad.Tensor.schedule_linear
 ::: tinygrad.Tensor.realize
 ::: tinygrad.Tensor.replace
 ::: tinygrad.Tensor.assign
@@ -0,0 +1,61 @@
+# TinyGPU
+
+TinyGPU app lets you use AMD and NVIDIA GPUs on macOS over USB4/Thunderbolt with tinygrad.
+
+## Requirements
+
+- macOS (13.0+)
+- USB4/Thunderbolt port
+- A supported GPU (AMD RDNA3+ or NVIDIA Ampere+)
+
+## Setup
+
+### 1. Connect your GPU
+
+Plug the supported GPU into your Mac over USB4/Thunderbolt.
+
+### 2. Initiate the driver install
+
+> **Note:** If tinygrad is cloned but not installed, run commands with `PYTHONPATH=.`
+
+```bash
+curl -fsSL https://raw.githubusercontent.com/tinygrad/tinygrad/master/extra/setup_tinygpu_osx.sh | sh
+```
+
+This downloads TinyGPU.app and triggers a system prompt to install the driver extension.
+
+### 3. Enable the driver
+
+You should see a system prompt: **"TinyGPU" would like to use a new driver extension**. Click **Open System Settings** and toggle TinyGPU on.
+
+If you missed the prompt, go to **System Settings > General > Login Items & Extensions > Driver Extensions** and toggle TinyGPU on.
+
+### 4. Compiler Setup
+
+#### AMD
+
+```bash
+curl -fsSL https://raw.githubusercontent.com/tinygrad/tinygrad/master/extra/setup_hipcomgr_osx.sh | sh
+```
+
+#### NV
+
+Install [Docker Desktop](https://www.docker.com/products/docker-desktop/) if you don't have it.
+
+```bash
+curl -fsSL https://raw.githubusercontent.com/tinygrad/tinygrad/master/extra/setup_nvcc_osx.sh | sh
+```
+
+Make sure `~/.local/bin` is on your `PATH`:
+
+```bash
+export PATH="$HOME/.local/bin:$PATH"
+```
+
+### 5. Use it!
+
+```bash
+DEV={AMD|NV} python3 -m tinygrad.llm
+```
+
+**Note:** Use `JITBEAM=2` to search for faster kernels (one-time search cost, results cached).
@@ -113,7 +113,7 @@ class VLIWRenderer(Renderer):
        case Ops.GEP:
          # a GEP is just an alias to a special register in the vector
          r[u] = r[u.src[0]] + u.arg[0]
-        case Ops.VECTORIZE:
+        case Ops.STACK:
          if all(s == u.src[0] for s in u.src):
            # if all sources are the same, we can broadcast
            inst.append({"valu": [("vbroadcast", r[u], r[u.src[0]])]})
@@ -173,16 +173,16 @@ if __name__ == "__main__":

  # *** render to device ***

-  from tinygrad.codegen import get_program
+  from tinygrad.codegen import to_program
  with Context(PCONTIG=2, DEVECTORIZE=2, SPEC=0):
    out = tree_traversal(forest_t, val_t, height, rounds)
-    sink = out.schedule()[-1].ast
-    prg = get_program(sink, VLIWRenderer())
+    sink = out.schedule_linear().src[-1].src[0]
+    prg = to_program(sink, VLIWRenderer())

  # *** run on Machine and compare ***

  # NOTE: the scratch size needs to be reduced to 1536 when you have a register allocator
-  src = eval(prg.src)
+  src = eval(prg.src[3].arg)
  max_regs = max(t[1] for instr in src for v in instr.values() for t in v if len(t) > 1) + 8
  print(f"{max_regs:5d} regs used" + ("" if max_regs <= 1536 else "       <-- WARNING: TOO MANY REGISTERS, MUST BE <= 1536"))
  machine = problem.Machine(mem, src, problem.DebugInfo(scratch_map={}), n_cores=1, trace=False, scratch_size=max_regs)
@@ -35,12 +35,11 @@ def compile_onnx_model(onnx_model):
  tinyonnx = TinyOnnx(onnx_model)
  the_input = Tensor.randn(1,32)

-  run, special_names = jit_model(tinyonnx, the_input)
+  linear, output_bufs = jit_model(tinyonnx, the_input)
+  the_output = [tinyonnx.forward(the_input)]

-  functions, statements, bufs, bufs_to_save = compile_net(run, special_names)
+  functions, statements, bufs, bufs_to_save = compile_net(linear, output_bufs)
  prg = export_model_clang(functions, statements, bufs, {}, ["input0"], ["output0"])
-
-  the_output = run(the_input)
  cprog = ["#include <string.h>", "#include <stdio.h>", "#include <stdlib.h>"]
  cprog.append(prg)

@@ -5,8 +5,9 @@ with contextlib.suppress(ImportError): import tiktoken
 from tinygrad import Tensor, TinyJit, Device, GlobalCounters, Variable, dtypes
 from tinygrad.uop.ops import UOp
 from tinygrad.helpers import Timing, DEBUG, JIT, getenv, fetch, colored, trange
+from tinygrad.llm.gguf import gguf_load
 from tinygrad.nn import Embedding, Linear, LayerNorm
-from tinygrad.nn.state import gguf_load, torch_load, load_state_dict, get_state_dict
+from tinygrad.nn.state import torch_load, load_state_dict, get_state_dict
 from extra.bench_log import BenchEvent, WallTimeEvent

 MAX_CONTEXT = getenv("MAX_CONTEXT", 128)
@@ -445,7 +445,7 @@ After you are done speaking, output [EOS]. You are not Chad.
  print(f"using LLaMA{LLAMA_SUFFIX}-{args.size} model")
  device = tuple(f"{Device.DEFAULT}:{i}" for i in range(args.shard)) if args.shard > 1 else Device.DEFAULT
  llama = LLaMa.build(MODEL_PATH, TOKENIZER_PATH, model_gen=args.gen, model_size=args.size, quantize=args.quantize, device=device)
-  param_bytes = sum(x.uop.size * x.dtype.itemsize for x in get_parameters(llama.model))
+  param_bytes = sum(x.nbytes() for x in get_parameters(llama.model))

  outputted = pre_prompt if chatbot else args.prompt
  start_pos, toks = 0, [llama.tokenizer.bos_id()] + llama.tokenizer.encode(outputted)
@@ -2,7 +2,8 @@ from pathlib import Path
 from typing import List
 import json, argparse, random, time, os
 from extra.models.llama import Transformer, convert_from_huggingface, convert_from_gguf, fix_bf16
-from tinygrad.nn.state import safe_load, torch_load, load_state_dict, get_parameters, gguf_load
+from tinygrad.llm.gguf import gguf_load
+from tinygrad.nn.state import safe_load, torch_load, load_state_dict, get_parameters
 from tinygrad import Tensor, dtypes, nn, Context, Device, GlobalCounters
 from tinygrad.helpers import Profiling, Timing, DEBUG, colored, fetch, tqdm
 from extra.bench_log import BenchEvent, WallTimeEvent
@@ -122,7 +123,7 @@ def NF4Linear(block_size):
    def __call__(self, x: Tensor) -> Tensor:
      high_bits = self.weight
      low_bits = (self.weight * 2 ** 4).contiguous()
-      unpacked = Tensor.stack(high_bits, low_bits, dim=-1).idiv(2 ** 4)
+      unpacked = Tensor.stack(high_bits, low_bits, dim=-1).div(2 ** 4, rounding_mode="trunc")
      unscaled = CODE[unpacked].to(x.device).reshape(-1, block_size) * self.scale
      return x.linear(unscaled.reshape(self.out_features, self.in_features).T)

@@ -324,7 +325,7 @@ if __name__ == "__main__":

  device = tuple(f"{Device.DEFAULT}:{i}" for i in range(args.shard)) if args.shard > 1 else Device.DEFAULT
  model = build_transformer(args.model, model_size=args.size, quantize=args.quantize, device=device)
-  param_bytes = sum(x.uop.size * x.dtype.itemsize for x in get_parameters(model))
+  param_bytes = sum(x.nbytes() for x in get_parameters(model))

  if not args.no_api and not args.benchmark:
    from bottle import Bottle, request, response, HTTPResponse, abort, static_file
@@ -2,13 +2,14 @@
 import os
 if "NOOPT" not in os.environ: os.environ["NOOPT"] = "1"
 from tinygrad import Device, nn, Tensor, dtypes
-Device.DEFAULT = "CPU"
 from train_gpt2 import GPT, GPTConfig
-from tinygrad.helpers import dedup, flatten, getenv, GlobalCounters, to_function_name
+from tinygrad.helpers import DEV, dedup, flatten, getenv, GlobalCounters, to_function_name
 from tinygrad.engine.realize import get_kernel
-from tinygrad.engine.memory import memory_planner
+from tinygrad.schedule.memory import memory_planner
 from tinygrad.uop.ops import Ops

+DEV.value = "CPU"
+
 TIMING = getenv("TIMING")

 if __name__ == "__main__":
@@ -325,19 +325,18 @@ def eval_stable_diffusion():
  # NOTE: the clip weights are the same between model.cond_stage_model and clip_encoder
  eval_timesteps = list(reversed(range(1, 1000, 20)))

-  original_device, Device.DEFAULT = Device.DEFAULT, "CPU"
-  # The choice of alphas_prev[0] = alphas_cumprod[0] seems arbitrary, but it's how the mlperf ref does it:
-  #   alphas_prev = np.asarray([alphacums[0]] + alphacums[ddim_timesteps[:-1]].tolist())
-  eval_alphas_prev = model.alphas_cumprod[0:1].cat(model.alphas_cumprod[list(range(1, 1000, 20))[:-1]]).to(GPUS).realize()
-  inception = FidInceptionV3().load_from_pretrained(CKPTDIR / "inception" / "pt_inception-2015-12-05-6726825d.pth")
-  vision_cfg = {'width': 1280, 'layers': 32, 'd_head': 80, 'image_size': 224, 'patch_size': 14}
-  text_cfg = {'width': 1024, 'n_heads': 16, 'layers': 24, 'vocab_size': 49408, 'ctx_length': 77}
-  clip.gelu = gelu_erf
-  clip_encoder = OpenClipEncoder(1024, text_cfg, vision_cfg)
-  loaded = torch_load(CKPTDIR / "clip" / "open_clip_pytorch_model.bin")
-  loaded.update({"attn_mask": clip_encoder.attn_mask, "mean": clip_encoder.mean, "std": clip_encoder.std})
-  load_state_dict(clip_encoder, loaded)
-  Device.DEFAULT=original_device
+  with Context(DEV="CPU"):
+    # The choice of alphas_prev[0] = alphas_cumprod[0] seems arbitrary, but it's how the mlperf ref does it:
+    #   alphas_prev = np.asarray([alphacums[0]] + alphacums[ddim_timesteps[:-1]].tolist())
+    eval_alphas_prev = model.alphas_cumprod[0:1].cat(model.alphas_cumprod[list(range(1, 1000, 20))[:-1]]).to(GPUS).realize()
+    inception = FidInceptionV3().load_from_pretrained(CKPTDIR / "inception" / "pt_inception-2015-12-05-6726825d.pth")
+    vision_cfg = {'width': 1280, 'layers': 32, 'd_head': 80, 'image_size': 224, 'patch_size': 14}
+    text_cfg = {'width': 1024, 'n_heads': 16, 'layers': 24, 'vocab_size': 49408, 'ctx_length': 77}
+    clip.gelu = gelu_erf
+    clip_encoder = OpenClipEncoder(1024, text_cfg, vision_cfg)
+    loaded = torch_load(CKPTDIR / "clip" / "open_clip_pytorch_model.bin")
+    loaded.update({"attn_mask": clip_encoder.attn_mask, "mean": clip_encoder.mean, "std": clip_encoder.std})
+    load_state_dict(clip_encoder, loaded)

  @TinyJit
  def denoise_step(x:Tensor, x_x:Tensor, t_t:Tensor, uc_c:Tensor, sqrt_alphas_cumprod_t:Tensor, sqrt_one_minus_alphas_cumprod_t:Tensor,
@@ -246,7 +246,7 @@ def train_resnet():

      if i == BENCHMARK:
        assert not math.isnan(loss)
-        median_step_time = sorted(step_times)[(BENCHMARK + 1) // 2]  # in seconds
+        median_step_time = sorted(step_times)[BENCHMARK // 2]  # in seconds
        estimated_total_minutes = int(median_step_time * steps_in_train_epoch * epochs / 60)
        print(f"Estimated training time: {estimated_total_minutes // 60}h{estimated_total_minutes % 60}m")
        print(f"epoch global_ops: {steps_in_train_epoch * GlobalCounters.global_ops:_}, "
@@ -593,7 +593,7 @@ def train_retinanet():

      if i == BENCHMARK:
        assert not math.isnan(loss)
-        median_step_time = sorted(step_times)[(BENCHMARK + 1) // 2]  # in seconds
+        median_step_time = sorted(step_times)[BENCHMARK // 2]  # in seconds
        estimated_total_minutes = int(median_step_time * steps_in_train_epoch * EPOCHS / 60)
        print(f"Estimated training time: {estimated_total_minutes // 60}h{estimated_total_minutes % 60}m")
        print(f"epoch global_ops: {steps_in_train_epoch * GlobalCounters.global_ops:_}, "
@@ -868,7 +868,7 @@ def train_unet3d():
        i += 1

        if i == BENCHMARK:
-          median_step_time = sorted(step_times)[(BENCHMARK + 1) // 2]  # in seconds
+          median_step_time = sorted(step_times)[BENCHMARK // 2]  # in seconds
          estimated_total_minutes = int(median_step_time * SAMPLES_PER_EPOCH * NUM_EPOCHS / 60)
          print(f"Estimated training time: {estimated_total_minutes // 60}h{estimated_total_minutes % 60}m")
          if (TRAIN_BEAM or EVAL_BEAM) and epoch == start_epoch: break
@@ -1167,7 +1167,7 @@ def train_bert():
      i += 1

      if i == BENCHMARK:
-        median_step_time = sorted(step_times)[(BENCHMARK + 1) // 2]  # in seconds
+        median_step_time = sorted(step_times)[BENCHMARK // 2]  # in seconds
        estimated_total_minutes = int(median_step_time * train_steps / 60)
        print(f"Estimated training time: {estimated_total_minutes // 60}h{estimated_total_minutes % 60}m")
        print(f"epoch global_ops: {train_steps * GlobalCounters.global_ops:_}, "
@@ -1282,11 +1282,14 @@ def train_bert():
        previous_step = i

 def train_llama3():
-  from examples.mlperf.models.flat_llama import FlatTransformer
+  from examples.mlperf.models.flat_llama import FlatTransformer, apply_grad, FP8_DTYPE
  from examples.llama3 import MODEL_PARAMS
  from examples.mlperf.lr_schedulers import CosineAnnealingLRWithWarmup
  from examples.mlperf.optim import GradAccClipAdamW

+  INITMLPERF = getenv("INITMLPERF")
+  RUNMLPERF = getenv("RUNMLPERF")
+  LOGMLPERF = getenv("LOGMLPERF")
  BENCHMARK = getenv("BENCHMARK")

  config = {}
@@ -1309,15 +1312,61 @@ def train_llama3():
  EVAL_BS            = config["EVAL_BS"]                = getenv("EVAL_BS", 16)
  EVAL_TARGET        = config["EVAL_TARGET"]            = getenv("EVAL_TARGET", 5.6)

-  # LR=1e-4 TRAIN_ON_VAL=1 DEFAULT_FLOAT=bfloat16 JITBEAM=2 OPTIM_DTYPE=bfloat16 LLAMA3_SIZE=1B WARMUP_STEPS=36 DECAY_STEPS=360 SEQLEN=512 PYTHONPATH=. DEV=AMD AMD_LLVM=0 MODEL=llama3 python3 examples/mlperf/model_train.py
-  # trains to 7
+  if LOGMLPERF:
+    from mlperf_logging import mllog
+    import mlperf_logging.mllog.constants as mllog_constants
+
+    mllog.config(filename=f"result_llama31_{SEED}.log")
+    mllog.config(root_dir=Path(__file__).parents[3].as_posix())
+    MLLOGGER = mllog.get_mllogger()
+    MLLOGGER.logger.propagate = False
+
+    LLAMA_BENCHMARK = mllog_constants.LLAMA31_405B if getenv("LLAMA3_SIZE", "8B") == "405B" else mllog_constants.LLAMA31_8B
+
+    if INITMLPERF:
+      assert BENCHMARK, "BENCHMARK must be set for INITMLPERF"
+      MLLOGGER.event(key=mllog_constants.SUBMISSION_ORG, value="tinycorp")
+      MLLOGGER.event(key=mllog_constants.SUBMISSION_PLATFORM, value=getenv("SUBMISSION_PLATFORM", "tinybox"))
+      MLLOGGER.event(key=mllog_constants.SUBMISSION_DIVISION, value=mllog_constants.CLOSED)
+      MLLOGGER.event(key=mllog_constants.SUBMISSION_STATUS, value=mllog_constants.ONPREM)
+
+      MLLOGGER.event(key=mllog_constants.SUBMISSION_BENCHMARK, value=LLAMA_BENCHMARK)
+
+      diskcache_clear()
+      MLLOGGER.event(key=mllog_constants.CACHE_CLEAR, value=True)
+      MLLOGGER.start(key=mllog_constants.INIT_START, value=None)
+
+    if RUNMLPERF:
+      MLLOGGER.start(key=mllog_constants.RUN_START, value=None)
+      MLLOGGER.event(key=mllog_constants.SEED, value=SEED)
+
+      MLLOGGER.event(key=mllog_constants.GLOBAL_BATCH_SIZE, value=GBS)
+      MLLOGGER.event(key=mllog_constants.MAX_SEQUENCE_LENGTH, value=SEQLEN)
+      MLLOGGER.event(key=mllog_constants.MAX_STEPS, value=MAX_STEPS)
+      MLLOGGER.event(key=mllog_constants.GRADIENT_ACCUMULATION_STEPS, value=grad_acc)
+      MLLOGGER.event(key=mllog_constants.EVAL_SAMPLES, value=EVAL_SAMPLES)
+      MLLOGGER.event(key=mllog_constants.TRAIN_SAMPLES, value=SAMPLES)
+
+      MLLOGGER.event(key=mllog_constants.OPT_NAME, value=mllog_constants.ADAMW)
+      MLLOGGER.event(key=mllog_constants.OPT_BASE_LR, value=LR)
+      MLLOGGER.event(key=mllog_constants.OPT_END_LR, value=END_LR)
+      MLLOGGER.event(key=mllog_constants.OPT_ADAMW_BETA_1, value=0.9)
+      MLLOGGER.event(key=mllog_constants.OPT_ADAMW_BETA_2, value=0.95)
+      MLLOGGER.event(key=mllog_constants.OPT_ADAMW_EPSILON, value=1e-5)
+      MLLOGGER.event(key=mllog_constants.OPT_ADAMW_WEIGHT_DECAY, value=0.1)
+      MLLOGGER.event(key=mllog_constants.OPT_LR_WARMUP_STEPS, value=WARMUP_STEPS)
+      MLLOGGER.event(key=mllog_constants.NUM_WARMUP_STEPS, value=WARMUP_STEPS)
+      MLLOGGER.event(key=mllog_constants.OPT_LR_DECAY_STEPS, value=MAX_STEPS - WARMUP_STEPS)
+      MLLOGGER.event(key=mllog_constants.OPT_LR_DECAY_SCHEDULE, value="cosine with linear warmup")
+      MLLOGGER.event(key=mllog_constants.OPT_GRADIENT_CLIP_NORM, value=1.0)
+  else:
+    MLLOGGER = None

  opt_adamw_beta_1 = 0.9
  opt_adamw_beta_2 = 0.95
  opt_adamw_epsilon = 1e-5
  opt_adamw_weight_decay = 0.1

-  opt_gradient_clip_norm = 1.0
  opt_learning_rate_warmup_steps = WARMUP_STEPS
  opt_learning_rate_decay_steps = MAX_STEPS - opt_learning_rate_warmup_steps
  opt_base_learning_rate = LR
@@ -1347,9 +1396,9 @@ def train_llama3():

  params = get_parameters(model)

-  if getenv("FAKEDATA"):
+  if getenv("EMPTYWEIGHT"):
    for v in get_parameters(model):
-      v = v.assign(Tensor.empty(v.shape))
+      v = v.assign(Tensor.empty(v.shape, dtype=v.dtype))

  is_dp = (DP := getenv("DP", 1)) > 1
  is_mp = (MP := getenv("MP", 1)) > 1
@@ -1368,9 +1417,12 @@ def train_llama3():
  optim = GradAccClipAdamW(params, lr=0.0, b1=opt_adamw_beta_1, b2=opt_adamw_beta_2,
                           eps=opt_adamw_epsilon, weight_decay=opt_adamw_weight_decay, grad_acc=grad_acc, device=optim_device)

-  # init grads
  for p in optim.params:
-    p.grad = Tensor.zeros_like(p).contiguous()
+    grad_dtype = dtypes.bfloat16 if p.dtype == FP8_DTYPE else p.dtype
+    if isinstance(p.device, tuple) and p.uop.axis is not None:
+      p.grad = Tensor.zeros(p.shape, dtype=grad_dtype, device=p.device[0]).shard_(p.device, axis=p.uop.axis).contiguous()
+    else:
+      p.grad = Tensor.zeros(p.shape, dtype=grad_dtype, device=p.device).contiguous()
  grads = [p.grad for p in optim.params]

  scheduler = CosineAnnealingLRWithWarmup(optim, opt_base_learning_rate, opt_end_learning_rate, opt_learning_rate_warmup_steps, opt_learning_rate_decay_steps)
@@ -1384,19 +1436,40 @@ def train_llama3():
    print(f"loading optim checkpoint from {fn}")
    load_state_dict(scheduler, safe_load(fn), realize=False)

+  fp8_amax = [t for ts in model._fp8_amax.values() for t in ts]
+  fp8_grad_amax = [t for ts in model._fp8_grad_amax.values() for t in ts] if hasattr(model, "_fp8_grad_amax") else []
+  fp8_inv_scales = list(model._fp8_inv_scale.values())
+
+  from tinygrad.nn.state import get_state_dict
+  model_state = get_state_dict(model)
+  for wname in ["wqkv", "wo", "w13", "w2"]:
+    w = model_state[wname]
+    w._inv_scale = model._fp8_inv_scale[wname]
+    if optim.master_params:
+      idx = next(j for j, p in enumerate(optim.params) if p is w)
+      optim.master_params[idx].assign((optim.master_params[idx] * w._inv_scale.reshape(-1, *([1]*(w.ndim-1)))).contiguous())
+
+  # realize everything here
+  if optim.master_params: Tensor.realize(*optim.master_params)
+  Tensor.realize(*optim.params, *fp8_inv_scales, *fp8_amax, *fp8_grad_amax)
+
  @TinyJit
  def minibatch(tokens:Tensor):
    if is_dp: tokens = tokens.to(None).shard(device, 0)
    if is_mp: tokens = tokens.shard(device)
    if not is_sharding: tokens = tokens.to(None)
    logits:Tensor = model(tokens[:, :-1])
-    loss = vocab_mask.where(-1e9, logits).sparse_categorical_crossentropy(tokens[:, 1:])
+    if getenv("FAST_CE", 0):
+      from extra.llama_kernels.fused_ce import fused_ce_loss
+      loss = fused_ce_loss(logits.cast(dtypes.bfloat16), tokens[:, 1:], label_smoothing=0.0)
+    else:
+      loss = vocab_mask.where(-1e9, logits).sparse_categorical_crossentropy(tokens[:, 1:])

-    loss.backward()
-    assert all(p.grad is g for p,g in zip(optim.params, grads))
+    for g, new_g in zip(grads, loss.gradient(*optim.params)):
+      apply_grad(g, new_g.uop)

    loss_cpu = loss.flatten().float().to("CPU")
-    return loss_cpu.realize(*grads)
+    return loss_cpu.realize(*grads, *fp8_amax, *fp8_grad_amax)

  @TinyJit
  def optim_step():
@@ -1407,7 +1480,7 @@ def train_llama3():

    lr_cpu = optim.lr.float().to("CPU")
    grad_norm_cpu = grad_norm.float().to("CPU")
-    Tensor.realize(lr_cpu, grad_norm_cpu, *grads)
+    Tensor.realize(lr_cpu, grad_norm_cpu, *grads, *fp8_inv_scales)

    return lr_cpu, grad_norm_cpu

@@ -1451,6 +1524,11 @@ def train_llama3():
  train_iter = get_train_iter()
  i, sequences_seen = resume_ckpt, 0
  step_times = []
+
+  if MLLOGGER and RUNMLPERF:
+    MLLOGGER.start(key=mllog_constants.EPOCH_START, metadata={mllog_constants.SAMPLES_COUNT: sequences_seen})
+    MLLOGGER.start(key=mllog_constants.BLOCK_START, metadata={mllog_constants.SAMPLES_COUNT: sequences_seen})
+
  while i < MAX_STEPS:
    GlobalCounters.reset()
    actual_gbs = GBS if i >= 2 else BS
@@ -1489,7 +1567,7 @@ def train_llama3():

      mem_gb = GlobalCounters.mem_used / 1e9
      gflops = GlobalCounters.global_ops / 1e9 / dev_time
-      mfu = ((6 * num_params * SEQLEN * GBS) / (dev_time * device_count * 2.3e15)) * 100
+      mfu = ((6 * num_params * SEQLEN * GBS) / (dev_time * device_count * 4.6e15)) * 100
      tqdm.write(
          f"{i:5} {step_time:.3f} s step, {gbs_time:.3f} s gbs, {optim_time:.3f} s optim, {data_time:.3f} s data, {loss:.4f} loss, " \
          f"{lr:.12f} LR, {grad_norm:.6f} grad_norm, {mem_gb:.2f} GB used, {gflops:9.2f} GFLOPS, {mfu:5.2f}% MFU")
@@ -1522,8 +1600,9 @@ def train_llama3():
        safe_save(get_state_dict(scheduler), fn)

      if i == BENCHMARK:
-        median_step_time = sorted(step_times)[(BENCHMARK + 1) // 2]
-        estimated_total_minutes = int(median_step_time * (SAMPLES // GBS) / 60)
+        median_step_time = sorted(step_times)[BENCHMARK // 2]
+        estimated_steps = 200_000 // GBS if getenv("LLAMA3_SIZE", "8B") == "8B" else MAX_STEPS
+        estimated_total_minutes = int(median_step_time * estimated_steps / 60)
        print(f"Estimated training time: {estimated_total_minutes // 60}h{estimated_total_minutes % 60}m")
        print(f"epoch global_ops: {GlobalCounters.global_ops:_}, "
              f"epoch global_mem: {GlobalCounters.global_mem:_}")
@@ -1533,6 +1612,10 @@ def train_llama3():
      tqdm.write(f"evaluating after {sequences_seen} sequences")
      profile_marker(f"eval @ {i}")

+      if MLLOGGER and RUNMLPERF:
+        MLLOGGER.end(key=mllog_constants.BLOCK_STOP, metadata={mllog_constants.SAMPLES_COUNT: sequences_seen})
+        MLLOGGER.start(key=mllog_constants.EVAL_START, metadata={mllog_constants.SAMPLES_COUNT: sequences_seen})
+
      # run eval
      eval_losses = []
      eval_iter = get_eval_iter()
@@ -1542,22 +1625,33 @@ def train_llama3():
        eval_losses += eval_step(tokens).tolist()

        if BENCHMARK and (j+1) == min(BENCHMARK, EVAL_SAMPLES//EVAL_BS):
+          if MLLOGGER and INITMLPERF:
+            MLLOGGER.end(key=mllog_constants.INIT_STOP, value=None)
          return

      log_perplexity = sum(eval_losses) / len(eval_losses)

      tqdm.write(f"eval log perplexity: {log_perplexity:.4f}")

+      if MLLOGGER and RUNMLPERF:
+        MLLOGGER.event(key=mllog_constants.EVAL_ACCURACY, value=log_perplexity, metadata={mllog_constants.SAMPLES_COUNT: sequences_seen})
+        MLLOGGER.end(key=mllog_constants.EVAL_STOP, metadata={mllog_constants.SAMPLES_COUNT: sequences_seen})
+
      if WANDB:
        wandb.log({"eval/log_perplexity": log_perplexity, "eval/sequences_seen": sequences_seen})

      if log_perplexity < EVAL_TARGET:
        tqdm.write(f"target achieved after {sequences_seen} sequences")
+        if MLLOGGER and RUNMLPERF:
+          MLLOGGER.end(key=mllog_constants.EPOCH_STOP, metadata={mllog_constants.SAMPLES_COUNT: sequences_seen})
+          MLLOGGER.end(key=mllog_constants.RUN_STOP, metadata={mllog_constants.STATUS: mllog_constants.SUCCESS})
        if getenv("CKPT"):
          if not os.path.exists(ckpt_dir := "./ckpts"): os.mkdir(ckpt_dir)
          fn = f"{ckpt_dir}/llama3.safe"
          safe_save(get_state_dict(model), fn)
        break
+      if MLLOGGER and RUNMLPERF:
+        MLLOGGER.start(key=mllog_constants.BLOCK_START, metadata={mllog_constants.SAMPLES_COUNT: sequences_seen})

 def train_stable_diffusion():
  from extra.models.unet import UNetModel
@@ -16,29 +16,78 @@ from tinygrad import Tensor, nn, function, getenv, dtypes, TinyJit
 from tinygrad.helpers import Timing, colored, GlobalCounters, profile_marker
 from tinygrad.uop.ops import Ops, UOp
 from extra.models.llama import apply_rotary_emb, precompute_freqs_cis
+from extra.llama_kernels.rmsnorm import rmsnorm
+from extra.llama_kernels import FP8_MAX, local_abs_max

-FP8 = getenv("FP8", 0)
-WQKV = getenv("WQKV", 0)
+ASM_GEMM = getenv("ASM_GEMM", 0)
+FUSED_INPUT_QUANTIZE = getenv("FUSED_INPUT_QUANTIZE", 0)
+FUSED_ADD_NORM_MUL_QUANTIZE = getenv("FUSED_ADD_NORM_MUL_QUANTIZE", 0)
+FUSED_SILU_W13 = getenv("FUSED_SILU_W13", 0)

 FP8_DTYPE = dtypes.fp8e4m3
-FP8_MAX = 448.0
+FP8_GRAD_DTYPE = dtypes.fp8e5m2

-def quantize_fp8(x:Tensor):
-  scale = FP8_MAX / (x.abs().max().detach() + 1e-8)
+def quantize_fp8(x:Tensor, amax_state:Tensor|None=None):
+  new_amax = (local_abs_max(x) if isinstance(x.device, tuple) else x.abs().max()).detach().cast(dtypes.float32)
+  scale = FP8_MAX / ((amax_state if amax_state is not None else new_amax) + 1e-8)
  x_scaled = x * scale
  x_clamped = x_scaled + (x_scaled.detach().clamp(-FP8_MAX, FP8_MAX) - x_scaled.detach())  # STE
-  return x_clamped.cast(FP8_DTYPE), scale.float().reciprocal()
+  return x_clamped.cast(FP8_DTYPE), scale.float().reciprocal(), new_amax

-def matmul(x:Tensor, w:Tensor) -> Tensor:
-  if not FP8: return x @ w.T
-  # weights are already FP8, just quantize activations
-  x_fp8, x_scale = quantize_fp8(x)
-  return x_fp8.dot(w.T, dtype=dtypes.float) * x_scale
+def matmul(x:Tensor, w:Tensor, fp8:bool=True, amax_x:Tensor|None=None, w_inv_scale:Tensor|None=None,
+           x_fp8:Tensor|None=None, x_scale:Tensor|None=None, x_new_amax:Tensor|None=None,
+           grad_amax_state:Tensor|None=None) -> tuple[Tensor,...]:
+  if not fp8:
+    if ASM_GEMM:
+      from extra.gemm.cdna_asm_gemm import can_use_asm_gemm, asm_gemm
+      if can_use_asm_gemm(x, w.T): return (asm_gemm(x, w.T),)
+    return (x @ w.T,)
+  assert w_inv_scale is not None, "fp8 matmul requires w_inv_scale (weights must be stored in fp8 with per-tensor scale)"
+  if x_fp8 is None:
+    if FUSED_INPUT_QUANTIZE and amax_x is not None:
+      from extra.llama_kernels.quantize_fp8_delayed import quantize_fp8_delayed
+      x_fp8, x_scale, x_new_amax, _ = quantize_fp8_delayed(x, amax_x, FP8_DTYPE)
+    else:
+      x_fp8, x_scale, x_new_amax = quantize_fp8(x, amax_state=amax_x)
+  if ASM_GEMM:
+    from extra.gemm.cdna_asm_gemm import can_use_asm_gemm, asm_gemm
+    if can_use_asm_gemm(x_fp8, w.T):
+      return asm_gemm(x_fp8, w.T, x_scale=x_scale, w_scale=w_inv_scale, grad_amax_state=grad_amax_state), x_new_amax, x_fp8, w
+  return x_fp8.dot(w.T, dtype=dtypes.float) * x_scale * w_inv_scale, x_new_amax, x_fp8, w

-def rmsnorm(x_in:Tensor, eps:float):
-  x = x_in.float()
-  x = x * (x.square().mean(-1, keepdim=True) + eps).rsqrt()
-  return x.cast(x_in.dtype)
+def norm_quantize_matmul(x:Tensor, norm:Tensor, w:Tensor, w_inv_scale:Tensor, eps:float, amax_x:Tensor, grad_amax_state:Tensor):
+  if FUSED_ADD_NORM_MUL_QUANTIZE:
+    from extra.llama_kernels.fused_rmsnorm_mul_quantize_fp8 import fused_rmsnorm_mul_quantize_fp8
+    x_fp8, x_inv_scale, new_amax, x_normed, rrms = fused_rmsnorm_mul_quantize_fp8(x, norm, amax_x, eps, FP8_DTYPE)
+    out, *ret = matmul(None, w, w_inv_scale=w_inv_scale, x_fp8=x_fp8, x_scale=x_inv_scale, x_new_amax=new_amax, grad_amax_state=grad_amax_state)
+    return out, x_normed, rrms, ret
+  x_normed, rrms = rmsnorm(x, eps)
+  out, *ret = matmul(x_normed * norm, w, amax_x=amax_x, w_inv_scale=w_inv_scale, grad_amax_state=grad_amax_state)
+  return out, x_normed, rrms, ret
+
+def add_norm_quantize_matmul(x:Tensor, residual:Tensor, norm:Tensor, w:Tensor, w_inv_scale:Tensor, eps:float, amax_x:Tensor):
+  if FUSED_ADD_NORM_MUL_QUANTIZE:
+    from extra.llama_kernels.fused_rmsnorm_mul_quantize_fp8 import fused_add_rmsnorm_mul_quantize_fp8
+    x_fp8, x_inv_scale, new_amax, h, x_normed, rrms = fused_add_rmsnorm_mul_quantize_fp8(x, residual, norm, amax_x, eps, FP8_DTYPE)
+    out, *ret = matmul(None, w, w_inv_scale=w_inv_scale, x_fp8=x_fp8, x_scale=x_inv_scale, x_new_amax=new_amax)
+    return out, h, x_normed, rrms, ret
+  h = x + residual
+  x_normed, rrms = rmsnorm(h, eps)
+  out, *ret = matmul(x_normed * norm, w, amax_x=amax_x, w_inv_scale=w_inv_scale)
+  return out, h, x_normed, rrms, ret
+
+def silu_w13_quantize_matmul(x_w13:Tensor, w2:Tensor, s_2:Tensor,
+                             amax_x2:Tensor,
+                             grad_amax_xw13:Tensor, grad_amax_xout:Tensor):
+  if FUSED_SILU_W13:
+    from extra.llama_kernels.cast_amax import fused_quantize_fp8_w13
+    x2_fp8, x2_inv_scale, new_amax_x2 = fused_quantize_fp8_w13(x_w13, amax_x2, FP8_DTYPE, grad_amax_state=grad_amax_xw13)
+    out, *ret = matmul(None, w2, w_inv_scale=s_2, x_fp8=x2_fp8, x_scale=x2_inv_scale, x_new_amax=new_amax_x2, grad_amax_state=grad_amax_xout)
+    return out, ret
+  hidden = x_w13.shape[-1] // 2
+  x_w1, x_w3 = x_w13[..., :hidden], x_w13[..., hidden:]
+  out, *ret = matmul(x_w1.silu() * x_w3, w2, amax_x=amax_x2, w_inv_scale=s_2, grad_amax_state=grad_amax_xout)
+  return out, ret

 class FlatTransformer:
  def __init__(self, dim:int, hidden_dim:int, n_heads:int, n_layers:int, norm_eps:float, vocab_size:int, n_kv_heads:int|None=None,
@@ -49,22 +98,18 @@ class FlatTransformer:
    self.n_kv_heads = n_kv_heads if n_kv_heads is not None else n_heads # n_kv_heads != n_heads implies MQA [arxiv/2307.09288, A.2.1]
    self.head_dim = dim // n_heads
    self.n_rep = self.n_heads // self.n_kv_heads
+    self.hidden_dim = hidden_dim

    scaled_std = 0.02 / math.sqrt(2 * n_layers)

    # Attention
-    if WQKV:
-      self.wqkv = self.lin_per_layer(dim, self.n_heads * self.head_dim + self.n_kv_heads * self.head_dim * 2)
-    else:
-      self.wq = self.lin_per_layer(dim, self.n_heads * self.head_dim)
-      self.wk = self.lin_per_layer(dim, self.n_kv_heads * self.head_dim)
-      self.wv = self.lin_per_layer(dim, self.n_kv_heads * self.head_dim)
+    self._init_inv_scales = []  # populated by lin_per_layer
+    self.wqkv = self.lin_per_layer(dim, self.n_heads * self.head_dim + self.n_kv_heads * self.head_dim * 2)
    self.wo = self.lin_per_layer(self.n_heads * self.head_dim, dim, std=scaled_std)

    # FeedForward
-    self.w1 = self.lin_per_layer(dim, hidden_dim)
+    self.w13 = self.lin_per_layer(dim, hidden_dim * 2)
    self.w2 = self.lin_per_layer(hidden_dim, dim, std=scaled_std)
-    self.w3 = self.lin_per_layer(dim, hidden_dim)

    self.norm_eps = norm_eps
    self.attention_norm = Tensor.ones(n_layers, dim).contiguous()
@@ -73,51 +118,95 @@ class FlatTransformer:
    # output
    self.norm = nn.RMSNorm(dim, norm_eps)
    self.tok_embeddings = nn.Embedding(vocab_size, dim)
-    self.tok_embeddings.weight = Tensor.normal(vocab_size, dim, mean=0.0, std=0.02)
-    self.output = Tensor.normal(1, vocab_size, dim, mean=0.0, std=0.02)
+    self.tok_embeddings.weight = Tensor.normal(vocab_size, dim, mean=0.0, std=0.02, dtype=dtypes.bfloat16)
+    self.output = Tensor.normal(1, vocab_size, dim, mean=0.0, std=0.02, dtype=dtypes.bfloat16)
    self.freqs_cis = precompute_freqs_cis(dim // n_heads, max_context * 2, rope_theta).contiguous().requires_grad_(False)

+    def _amax(): return Tensor.full((), FP8_MAX, dtype=dtypes.float32).contiguous().requires_grad_(False)
+    names = ["xqkv", "xo", "x13", "x2"]
+    self._fp8_amax = {name: [_amax() for _ in range(n_layers)] for name in names}
+    grad_names = ["xqkv", "xo", "xw13", "xout"]
+    self._fp8_grad_amax = {name: [_amax() for _ in range(n_layers)] for name in grad_names}
+    w_names = ["wqkv", "wo", "w13", "w2"]
+    self._fp8_inv_scale = {wname: inv_scales.float().contiguous().requires_grad_(False)
+                           for wname, inv_scales in zip(w_names, self._init_inv_scales)}
+    del self._init_inv_scales
+
  def lin_per_layer(self, in_features:int, out_features:int, std:float=0.02):
-    dt = FP8_DTYPE if FP8 else None
-    if getenv("ZEROS"): return Tensor.zeros(self.n_layers, out_features, in_features, dtype=dt)
-    return Tensor.normal(self.n_layers, out_features, in_features, mean=0.0, std=std, dtype=dt)
+    if getenv("ZEROS"): w = Tensor.zeros(self.n_layers, out_features, in_features)
+    else: w = Tensor.normal(self.n_layers, out_features, in_features, mean=0.0, std=std)
+    amax = w.abs().flatten(1).max(1).detach()
+    scale = FP8_MAX / (amax + 1e-8)
+    self._init_inv_scales.append((amax + 1e-8) / FP8_MAX)
+    return (w * scale.reshape(-1, 1, 1)).clamp(-FP8_MAX, FP8_MAX).cast(FP8_DTYPE)

-  def attention(self, x:Tensor, freqs_cis:Tensor, attention_norm:Tensor, wo:Tensor, wqkv:Tensor|None=None,
-                wq:Tensor|None=None, wk:Tensor|None=None, wv:Tensor|None=None):
-    x = rmsnorm(x, self.norm_eps) * attention_norm
+  def attention(self, x:Tensor, freqs_cis:Tensor, attention_norm:Tensor, wqkv:Tensor, wo:Tensor,
+                amax_xqkv:Tensor, amax_xo:Tensor, s_qkv:Tensor, s_o:Tensor,
+                grad_amax_xqkv:Tensor, grad_amax_xo:Tensor):
    bsz, seqlen, _ = x.shape
+    new_amaxs, saves = [], []

-    if wqkv is not None:
-      xqkv = matmul(x, wqkv)
-      xqkv = xqkv.reshape(bsz, seqlen, self.n_kv_heads, self.n_rep + 2, self.head_dim)
-      xq = xqkv[:, :, :, :self.n_rep].reshape(bsz, seqlen, self.n_heads, self.head_dim)
-      xk = xqkv[:, :, :, self.n_rep].reshape(bsz, seqlen, self.n_kv_heads, self.head_dim)
-      xv = xqkv[:, :, :, self.n_rep+1].reshape(bsz, seqlen, self.n_kv_heads, self.head_dim)
-    else:
-      assert wq is not None and wk is not None and wv is not None
-      xq = matmul(x, wq).reshape(bsz, seqlen, self.n_heads, self.head_dim)
-      xk = matmul(x, wk).reshape(bsz, seqlen, self.n_kv_heads, self.head_dim)
-      xv = matmul(x, wv).reshape(bsz, seqlen, self.n_kv_heads, self.head_dim)
+    xqkv, x_normed, rrms, ret = norm_quantize_matmul(x, attention_norm, wqkv, s_qkv, self.norm_eps,
+                                                     amax_x=amax_xqkv, grad_amax_state=grad_amax_xqkv)
+    saves.extend([x_normed, rrms])
+    new_amaxs.extend(ret[:1])
+    saves.extend(ret[1:] + [xqkv])
+    xqkv = xqkv.reshape(bsz, seqlen, self.n_kv_heads, self.n_rep + 2, self.head_dim)
+    xq = xqkv[:, :, :, :self.n_rep].reshape(bsz, seqlen, self.n_heads, self.head_dim)
+    xk = xqkv[:, :, :, self.n_rep].reshape(bsz, seqlen, self.n_kv_heads, self.head_dim)
+    xv = xqkv[:, :, :, self.n_rep+1].reshape(bsz, seqlen, self.n_kv_heads, self.head_dim)

    xq, xk = apply_rotary_emb(xq, xk, freqs_cis)
-    xq, xk, xv = xq.transpose(1, 2), xk.transpose(1, 2), xv.transpose(1, 2)
-    attn = xq.scaled_dot_product_attention(xk, xv, is_causal=True, enable_gqa=True).transpose(1, 2)
+    xq, xk, xv = xq.cast(dtypes.bfloat16), xk.cast(dtypes.bfloat16), xv.cast(dtypes.bfloat16)
+    if getenv("HK_FLASH_ATTENTION"):
+      from extra.thunder.amd.fa import flash_attention
+      attn, *save = flash_attention(xq, xk, xv, is_causal=True)
+      saves.extend(save)
+    else:
+      xq, xk, xv = xq.transpose(1, 2), xk.transpose(1, 2), xv.transpose(1, 2)
+      attn = xq.scaled_dot_product_attention(xk, xv, is_causal=True, enable_gqa=True).transpose(1, 2)
    attn = attn.reshape(bsz, seqlen, -1)
-    return matmul(attn, wo)

-  def feed_forward(self, x:Tensor, ffn_norm:Tensor, w1:Tensor, w2:Tensor, w3:Tensor):
-    x = rmsnorm(x, self.norm_eps) * ffn_norm
-    x_w1 = matmul(x, w1).silu()
-    x_w3 = matmul(x.contiguous_backward(), w3)
-    return matmul(x_w1 * x_w3, w2)
+    out, *ret = matmul(attn, wo, amax_x=amax_xo, w_inv_scale=s_o, grad_amax_state=grad_amax_xo)
+    new_amaxs.extend(ret[:1])
+    saves.extend(ret[1:] + [out])
+    return (out, *new_amaxs, *saves)
+
+  def feed_forward(self, x:Tensor, residual:Tensor, ffn_norm:Tensor, w13:Tensor, w2:Tensor,
+                   amax_x13:Tensor, amax_x2:Tensor, s_13:Tensor, s_2:Tensor,
+                   grad_amax_xw13:Tensor, grad_amax_xout:Tensor):
+    new_amaxs, saves = [], []
+
+    x_w13, h, x_normed, rrms, ret = add_norm_quantize_matmul(x, residual, ffn_norm, w13, s_13, self.norm_eps,
+                                                             amax_x=amax_x13)
+    saves.extend([x_normed, rrms])
+    new_amaxs.extend(ret[:1])
+    saves.extend(ret[1:] + [x_w13])
+
+    out, ret = silu_w13_quantize_matmul(x_w13, w2, s_2, amax_x2=amax_x2, grad_amax_xw13=grad_amax_xw13, grad_amax_xout=grad_amax_xout)
+    new_amaxs.extend(ret[:1])
+    saves.extend(ret[1:] + [out])
+    return (out, h, *new_amaxs, *saves)

  @function(precompile=True, precompile_backward=True)
  def run_layer(self, x:Tensor, freqs_cis:Tensor,
-                attention_norm:Tensor, wo:Tensor,
-                ffn_norm:Tensor, w1:Tensor, w2:Tensor, w3:Tensor,
-                wqkv:Tensor|None=None, wq:Tensor|None=None, wk:Tensor|None=None, wv:Tensor|None=None):
-    h = x + self.attention(x, freqs_cis, attention_norm, wo, wqkv=wqkv, wq=wq, wk=wk, wv=wv)
-    return h + self.feed_forward(h, ffn_norm, w1, w2, w3)
+                attention_norm:Tensor, wqkv:Tensor, wo:Tensor,
+                ffn_norm:Tensor, w13:Tensor, w2:Tensor,
+                amax_xqkv:Tensor, amax_xo:Tensor,
+                amax_x13:Tensor, amax_x2:Tensor,
+                s_qkv:Tensor, s_o:Tensor, s_13:Tensor, s_2:Tensor,
+                grad_amax_xqkv:Tensor, grad_amax_xo:Tensor,
+                grad_amax_xw13:Tensor, grad_amax_xout:Tensor):
+    attn, *attn_ret = self.attention(x, freqs_cis, attention_norm, wqkv, wo,
+                                     amax_xqkv=amax_xqkv, amax_xo=amax_xo, s_qkv=s_qkv, s_o=s_o,
+                                     grad_amax_xqkv=grad_amax_xqkv, grad_amax_xo=grad_amax_xo)
+    attn_amaxs, attn_saves = attn_ret[:2], attn_ret[2:]
+    ffn, h, *ffn_ret = self.feed_forward(x, attn, ffn_norm, w13, w2,
+                                                   amax_x13=amax_x13, amax_x2=amax_x2, s_13=s_13, s_2=s_2,
+                                                   grad_amax_xw13=grad_amax_xw13, grad_amax_xout=grad_amax_xout)
+    ffn_amaxs, ffn_saves = ffn_ret[:2], ffn_ret[2:]
+    h = h + ffn
+    return (h, *attn_amaxs, *ffn_amaxs, *attn_saves, *ffn_saves)

  def shard(self, device:tuple[str, ...], mp:bool=False):
    from tinygrad.nn.state import get_parameters
@@ -125,43 +214,63 @@ class FlatTransformer:
      for v in get_parameters(self): v.shard_(device, axis=None)
    else:
      # flat per-layer weights: axis 0 is n_layers, so shard axes are +1 vs per-layer Transformer
-      if WQKV:
-        self.wqkv.shard_(device, axis=1).realize()          # (n_layers, out, dim) shard out
-      else:
-        self.wq.shard_(device, axis=1).realize()            # (n_layers, n_heads*head_dim, dim) shard out
-        self.wk.shard_(device, axis=1).realize()            # (n_layers, n_kv_heads*head_dim, dim) shard out
-        self.wv.shard_(device, axis=1).realize()            # (n_layers, n_kv_heads*head_dim, dim) shard out
-      self.wo.shard_(device, axis=2).realize()             # (n_layers, dim, in) shard in
-      self.w1.shard_(device, axis=1).realize()             # (n_layers, hidden, dim) shard out
-      self.w2.shard_(device, axis=2).realize()             # (n_layers, dim, hidden) shard in
-      self.w3.shard_(device, axis=1).realize()             # (n_layers, hidden, dim) shard out
+      self.wqkv.shard_(device, axis=1).realize()          # (n_layers, out, dim) shard out
+      self.wo.shard_(device, axis=2).realize()            # (n_layers, dim, in) shard in
+      self.w13.shard_(device, axis=1).realize()           # (n_layers, hidden*2, dim) shard out
+      self.w2.shard_(device, axis=2).realize()            # (n_layers, dim, hidden) shard in
      self.attention_norm.shard_(device, axis=None).realize()
      self.ffn_norm.shard_(device, axis=None).realize()
      self.norm.weight.shard_(device, axis=None).realize()
      self.tok_embeddings.weight.shard_(device, axis=0).realize()
      self.output.shard_(device, axis=1).realize()
      self.freqs_cis.shard_(device, axis=None).realize()
+      for amax_dict in (self._fp8_amax, self._fp8_grad_amax):
+        for name in amax_dict:
+          for i in range(len(amax_dict[name])):
+            amax_dict[name][i] = amax_dict[name][i].to(device).contiguous().requires_grad_(False)
+      for name in self._fp8_inv_scale:
+        self._fp8_inv_scale[name] = self._fp8_inv_scale[name].to(device).contiguous().requires_grad_(False)

  def __call__(self, tokens:Tensor):
    h = self.tok_embeddings(tokens)
    freqs_cis = self.freqs_cis.cast(h.dtype)[:, :tokens.shape[1], :, :, :]
+    a, ga, s = self._fp8_amax, self._fp8_grad_amax, self._fp8_inv_scale
    for i in range(self.n_layers):
-      attn_kwargs = {"wqkv": self.wqkv[i]} if WQKV else {"wq": self.wq[i], "wk": self.wk[i], "wv": self.wv[i]}
-      h = self.run_layer(h, freqs_cis,
-                         self.attention_norm[i], self.wo[i],
-                         self.ffn_norm[i], self.w1[i], self.w2[i], self.w3[i], **attn_kwargs)
-    logits = self.norm(h) @ self.output[0].T
+      h, *ret = self.run_layer(h, freqs_cis,
+                               self.attention_norm[i], self.wqkv[i], self.wo[i],
+                               self.ffn_norm[i], self.w13[i], self.w2[i],
+                               amax_xqkv=a["xqkv"][i], amax_xo=a["xo"][i],
+                               amax_x13=a["x13"][i], amax_x2=a["x2"][i],
+                               s_qkv=s["wqkv"][i], s_o=s["wo"][i],
+                               s_13=s["w13"][i], s_2=s["w2"][i],
+                               grad_amax_xqkv=ga["xqkv"][i], grad_amax_xo=ga["xo"][i],
+                               grad_amax_xw13=ga["xw13"][i], grad_amax_xout=ga["xout"][i])
+      for name, new_val in zip(["xqkv", "xo", "x13", "x2"], ret[:5]):
+        a[name][i].assign(new_val)
+
+    logits = matmul(self.norm(h), self.output[0], fp8=False)[0]
    return logits

-# TODO: this shouldn't be needed, but it prevents a copy of the grads. CAT can help
-def apply_grad(old_grad:UOp, new_grad:UOp) -> list[UOp]:
-  if new_grad.op == Ops.ADD:
-    return apply_grad(old_grad, new_grad.src[0])+apply_grad(old_grad, new_grad.src[1])
-  elif new_grad.op == Ops.PAD:
-    grad_shrink = tuple([(p[0], s+p[0]) for s,p in zip(new_grad.src[0].shape, new_grad.marg)])
-    return apply_grad(old_grad.shrink(grad_shrink), new_grad.src[0])
-  else:
-    return [old_grad.store(old_grad + new_grad)]
+def _get_pads(uop:UOp) -> list[UOp]:
+  if uop.op == Ops.ADD: return _get_pads(uop.src[0]) + _get_pads(uop.src[1])
+  return [uop]
+
+def apply_grad(grad_buf:Tensor, new_grad:UOp):
+  pads = _get_pads(new_grad)
+  if len(pads) <= 1:
+    new_grad = new_grad.cast(grad_buf.dtype)
+    store = grad_buf.uop.store(grad_buf.uop + new_grad)
+    grad_buf.uop = grad_buf.uop.after(store)
+    return
+  sorted_pads = sorted(pads, key=lambda p: p.marg[0][0] if p.op == Ops.PAD else 0)
+  inners_raw = [Tensor(p.src[0] if p.op == Ops.PAD else p, device=grad_buf.device) for p in sorted_pads]
+  if getenv("FUSED_PAD_GRAD_ACCUM", 0):
+    from extra.llama_kernels.fused_pad_grad_accum import fused_pad_grad_accum, can_fused_pad_grad_accum
+    if can_fused_pad_grad_accum(grad_buf, inners_raw):
+      grad_buf.uop = fused_pad_grad_accum(grad_buf, inners_raw).uop
+      return
+  inners = [t.cast(grad_buf.dtype) for t in inners_raw]
+  grad_buf.assign(grad_buf + inners[0].cat(*inners[1:], dim=0))

 if __name__ == "__main__":
  config = {}
@@ -183,7 +292,8 @@ if __name__ == "__main__":
    model.shard(tuple(f"{Device.DEFAULT}:{i}" for i in range(MP)), mp=True)

  # preallocate all the grad buffers and zero them out
-  grads = {x:Tensor.zeros_like(x).contiguous() for x in state.values() if x.requires_grad is None}
+  grads = {x:Tensor.zeros(x.shape, dtype=x.dtype, device=x.device).contiguous()
+           for x in state.values() if x.requires_grad is None}

  # print model size
  sz = 0
@@ -203,7 +313,7 @@ if __name__ == "__main__":
    with Timing("python forward: "): loss = model(tokens[:, :-1]).sparse_categorical_crossentropy(tokens[:, 1:])
    with Timing("python backward: "):
      for t,g in zip(grads, loss.gradient(*grads)):
-        grads[t] = Tensor(grads[t].uop.after(UOp.group(*apply_grad(grads[t].uop, g.uop))), device=t.device)
+        apply_grad(grads[t], g.uop)
    with Timing("run step: "): loss.realize(*grads.values())

  for i in range(6):
@@ -1,80 +0,0 @@
-from tinygrad import Tensor, nn
-from tinygrad.helpers import getenv
-from extra.models.llama import apply_rotary_emb, precompute_freqs_cis
-
-class Attention:
-  def __init__(self, dim:int, n_heads:int, n_kv_heads:int|None=None, linear=nn.Linear):
-    self.n_heads = n_heads
-    self.n_kv_heads = n_kv_heads if n_kv_heads is not None else n_heads # n_kv_heads != n_heads implies MQA [arxiv/2307.09288, A.2.1]
-    self.head_dim = dim // n_heads
-    self.n_rep = self.n_heads // self.n_kv_heads
-
-    if getenv("WQKV"):
-      self.wqkv = linear(dim, self.n_heads * self.head_dim + self.n_kv_heads * self.head_dim * 2, bias=False)
-    else:
-      self.wq = linear(dim, self.n_heads * self.head_dim, bias=False)
-      self.wk = linear(dim, self.n_kv_heads * self.head_dim, bias=False)
-      self.wv = linear(dim, self.n_kv_heads * self.head_dim, bias=False)
-
-    self.wo = linear(self.n_heads * self.head_dim, dim, bias=False)
-
-  def __call__(self, x:Tensor, freqs_cis:Tensor) -> Tensor:
-    if getenv("WQKV"):
-      xqkv = self.wqkv(x)
-      xqkv = xqkv.reshape(xqkv.shape[0], xqkv.shape[1], self.n_kv_heads, self.n_rep + 2, self.head_dim)
-      xq = xqkv[:, :, :, :self.n_rep].reshape(xqkv.shape[0], xqkv.shape[1], -1)
-      xk = xqkv[:, :, :, self.n_rep:self.n_rep+1].reshape(xqkv.shape[0], xqkv.shape[1], -1)
-      xv = xqkv[:, :, :, self.n_rep+1:self.n_rep+2].reshape(xqkv.shape[0], xqkv.shape[1], -1)
-    else:
-      xq, xk, xv = self.wq(x), self.wk(x), self.wv(x)
-
-    xq = xq.reshape(xq.shape[0], xq.shape[1], self.n_heads, self.head_dim)
-    xk = xk.reshape(xk.shape[0], xk.shape[1], self.n_kv_heads, self.head_dim)
-    xv = xv.reshape(xv.shape[0], xv.shape[1], self.n_kv_heads, self.head_dim)
-
-    xq, xk = apply_rotary_emb(xq, xk, freqs_cis)
-    bsz, seqlen, _, _ = xq.shape
-
-    xq, xk, xv = xq.transpose(1, 2), xk.transpose(1, 2), xv.transpose(1, 2)
-    attn = xq.scaled_dot_product_attention(xk, xv, is_causal=True, enable_gqa=True).transpose(1, 2)
-
-    attn = attn.reshape(bsz, seqlen, -1)
-    return self.wo(attn)
-
-class FeedForward:
-  def __init__(self, dim:int, hidden_dim:int, linear=nn.Linear):
-    self.w1 = linear(dim, hidden_dim, bias=False)
-    self.w2 = linear(hidden_dim, dim, bias=False)
-    self.w3 = linear(dim, hidden_dim, bias=False) # the gate in Gated Linear Unit
-
-  def __call__(self, x:Tensor) -> Tensor:
-    w1 = self.w1(x).silu()
-    w3 = self.w3(x)
-    return self.w2(w1 * w3)
-
-class TransformerBlock:
-  def __init__(self, dim:int, hidden_dim:int, n_heads:int, n_kv_heads:int|None, norm_eps:float, linear=nn.Linear):
-    self.attention = Attention(dim, n_heads, n_kv_heads, linear)
-    self.feed_forward = FeedForward(dim, hidden_dim, linear)
-    self.attention_norm = nn.RMSNorm(dim, norm_eps)
-    self.ffn_norm = nn.RMSNorm(dim, norm_eps)
-
-  def __call__(self, x:Tensor, freqs_cis:Tensor):
-    h = x + self.attention(self.attention_norm(x), freqs_cis)
-    return h + self.feed_forward(self.ffn_norm(h))
-
-class Transformer:
-  def __init__(self, dim:int, hidden_dim:int, n_heads:int, n_layers:int, norm_eps:float, vocab_size:int, n_kv_heads:int|None=None,
-               rope_theta:int=10000, max_context:int=1024, linear=nn.Linear, embedding=nn.Embedding):
-    self.layers = [TransformerBlock(dim, hidden_dim, n_heads, n_kv_heads, norm_eps, linear) for _ in range(n_layers)]
-    self.norm = nn.RMSNorm(dim, norm_eps)
-    self.tok_embeddings = embedding(vocab_size, dim)
-    self.output = nn.Linear(dim, vocab_size, bias=False) if embedding == nn.Embedding else linear(dim, vocab_size, bias=False)
-    self.freqs_cis = precompute_freqs_cis(dim // n_heads, max_context * 2, rope_theta).contiguous().requires_grad_(False)
-
-  def __call__(self, tokens:Tensor):
-    h = self.tok_embeddings(tokens)
-    freqs_cis = self.freqs_cis.cast(h.dtype)[:, :tokens.shape[1], :, :, :]
-    for layer in self.layers: h = layer(h, freqs_cis)
-    logits = self.output(self.norm(h))
-    return logits
@@ -34,7 +34,9 @@ class GradAccClipAdamW(Optimizer):
    else:
      updates, extra = self._step([], grads)
    for i, tt in enumerate(self.params): tt.assign(self._apply_update(tt, updates[i], self.master_params[i] if self.master_params else None))
-    to_realize = extra+self.params+self.buffers+(self.master_params or [])
+    # collect inv_scale tensors attached to fp8 params (set by _apply_update)
+    fp8_inv_scales = [tt._inv_scale for tt in self.params if hasattr(tt, '_inv_scale')]
+    to_realize = extra+self.params+self.buffers+(self.master_params or [])+fp8_inv_scales

    Tensor.realize(*to_realize)
    return extra[-1]
@@ -77,4 +79,12 @@ class GradAccClipAdamW(Optimizer):
    new_w = w.detach() - up
    if master is not None: master.assign(new_w)
    if STOCHASTIC_ROUND and t.dtype == dtypes.bfloat16: return stochastic_round_bf16(new_w)
+    if t.dtype in dtypes.fp8s:
+      from examples.mlperf.models.flat_llama import FP8_MAX
+      amax = new_w.float().abs().max(axis=tuple(range(1, new_w.ndim))).detach()  # per-layer amax for (n_layers, out, in)
+      scale = FP8_MAX / (amax + 1e-8)
+      fp8_w = (new_w * scale.reshape(-1, *([1]*(new_w.ndim-1)))).clamp(-FP8_MAX, FP8_MAX).cast(t.dtype)
+      if hasattr(t, '_inv_scale'):
+        t._inv_scale.assign(((amax + 1e-8) / FP8_MAX).cast(t._inv_scale.dtype))
+      return fp8_w
    return new_w.cast(t.dtype)
@@ -2,7 +2,6 @@

 export PYTHONPATH="."
 export DEV=${DEV:-AMD}
-export EMULATE="AMD_CDNA4"
 export CHECK_OOB=0
 export REWRITE_STACK_LIMIT=5000000 HCQDEV_WAIT_TIMEOUT_MS=240000
 export DEVICE_IN_FUNCTION_BUG=1
@@ -10,14 +9,22 @@ export DEVICE_IN_FUNCTION_BUG=1
 export DEBUG=${DEBUG:-2}
 export HK_FLASH_ATTENTION=${HK_FLASH_ATTENTION:-1}
 export ALL2ALL=${ALL2ALL:-1}
-export USE_ATOMICS=${USE_ATOMICS:-0}
+export USE_ATOMICS=${USE_ATOMICS:-1}
 export ASM_GEMM=${ASM_GEMM:-1}
 export WQKV=${WQKV:-1}
+export MASTER_WEIGHTS=${MASTER_WEIGHTS:-1}
+export FP8=${FP8:-1}
+export ALLREDUCE_CAST=${ALLREDUCE_CAST:-1}
+export FAST_CE=${FAST_CE:-0}
+export FUSED_INPUT_QUANTIZE=${FUSED_INPUT_QUANTIZE:-1}
+export FUSED_ADD_NORM_MUL_QUANTIZE=${FUSED_ADD_NORM_MUL_QUANTIZE:-1}
+export FUSED_SILU_W13=${FUSED_SILU_W13:-1}
+export FUSED_PAD_GRAD_ACCUM=${FUSED_PAD_GRAD_ACCUM:-1}
 export OFFLOAD_OPTIM=${OFFLOAD_OPTIM:-1}

 export DEFAULT_FLOAT="bfloat16" OPTIM_DTYPE="bfloat16"
-export DP=${DP:-1} MP=${MP:-8}
-export BS=${BS:-1} EVAL_BS=${EVAL_BS:-1} GRADIENT_ACC_STEPS=${GRADIENT_ACC_STEPS:-2}
+export DP=${DP:-1} MP=${MP:-8} BS=${BS:-1} EVAL_BS=${EVAL_BS:-1} GRADIENT_ACC_STEPS=${GRADIENT_ACC_STEPS:-2}
+export GBS=$((BS * GRADIENT_ACC_STEPS))

 export MODEL="llama3"
 export BASEDIR="/raid/datasets/c4/"
@@ -30,7 +37,7 @@ export DATA_SEED=${DATA_SEED:-5760}
 export JITBEAM=${JITBEAM:-3}
 export BEAM_UOPS_MAX=6000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=1

-export FAKEDATA=1 BENCHMARK=10
+export FAKEDATA=${FAKEDATA:-1} BENCHMARK=${BENCHMARK:-10}
 if [ -z "$FULL_LAYERS" ]; then
  export LLAMA_LAYERS=2
 fi
@@ -2,7 +2,6 @@

 export PYTHONPATH="."
 export DEV=${DEV:-AMD}
-export EMULATE="AMD_CDNA4"
 export CHECK_OOB=0
 export REWRITE_STACK_LIMIT=5000000 HCQDEV_WAIT_TIMEOUT_MS=240000
 export DEVICE_IN_FUNCTION_BUG=1
@@ -10,12 +9,21 @@ export DEVICE_IN_FUNCTION_BUG=1
 export DEBUG=${DEBUG:-2}
 export HK_FLASH_ATTENTION=${HK_FLASH_ATTENTION:-1}
 export ALL2ALL=${ALL2ALL:-1}
+export LATE_ALLREDUCE=${LATE_ALLREDUCE:-0}
 export USE_ATOMICS=${USE_ATOMICS:-1}
 export ASM_GEMM=${ASM_GEMM:-1}
-export WQKV=${WQKV:-0}
+export WQKV=${WQKV:-1}
+export MASTER_WEIGHTS=${MASTER_WEIGHTS:-1}
+export FP8=${FP8:-1}
+export ALLREDUCE_CAST=${ALLREDUCE_CAST:-1}
+export FAST_CE=${FAST_CE:-1}
+export FUSED_INPUT_QUANTIZE=${FUSED_INPUT_QUANTIZE:-1}
+export FUSED_ADD_NORM_MUL_QUANTIZE=${FUSED_ADD_NORM_MUL_QUANTIZE:-1}
+export FUSED_SILU_W13=${FUSED_SILU_W13:-1}
+export FUSED_PAD_GRAD_ACCUM=${FUSED_PAD_GRAD_ACCUM:-1}

 export DEFAULT_FLOAT="bfloat16" OPTIM_DTYPE="bfloat16"
-export DP=${DP:-8} MP=${MP:-1} BS=${BS:-8} EVAL_BS=${EVAL_BS:-8} GRADIENT_ACC_STEPS=${GRADIENT_ACC_STEPS:-4}
+export DP=${DP:-8} MP=${MP:-1} BS=${BS:-16} EVAL_BS=${EVAL_BS:-8} GRADIENT_ACC_STEPS=${GRADIENT_ACC_STEPS:-2}
 export GBS=$((BS * GRADIENT_ACC_STEPS))

 export MODEL="llama3"
@@ -34,7 +42,7 @@ export DATA_SEED=${DATA_SEED:-5760}
 export JITBEAM=${JITBEAM:-3}
 export BEAM_UOPS_MAX=6000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=1

-export FAKEDATA=1 BENCHMARK=10
+export FAKEDATA=${FAKEDATA:-1} BENCHMARK=${BENCHMARK:-10}
 if [ -z "$FULL_LAYERS" ]; then
  export LLAMA_LAYERS=2
 fi
@@ -2,7 +2,6 @@

 export PYTHONPATH="."
 export DEV=${DEV:-AMD}
-export EMULATE="AMD_CDNA4"
 export CHECK_OOB=0
 export REWRITE_STACK_LIMIT=5000000 HCQDEV_WAIT_TIMEOUT_MS=240000
 export DEVICE_IN_FUNCTION_BUG=1
@@ -10,9 +9,17 @@ export DEVICE_IN_FUNCTION_BUG=1
 export DEBUG=${DEBUG:-2}
 export HK_FLASH_ATTENTION=${HK_FLASH_ATTENTION:-1}
 export ALL2ALL=${ALL2ALL:-1}
-export USE_ATOMICS=${USE_ATOMICS:-0}
+export USE_ATOMICS=${USE_ATOMICS:-1}
 export ASM_GEMM=${ASM_GEMM:-1}
 export WQKV=${WQKV:-1}
+export MASTER_WEIGHTS=${MASTER_WEIGHTS:-1}
+export FP8=${FP8:-1}
+export ALLREDUCE_CAST=${ALLREDUCE_CAST:-1}
+export FAST_CE=${FAST_CE:-0}
+export FUSED_INPUT_QUANTIZE=${FUSED_INPUT_QUANTIZE:-1}
+export FUSED_ADD_NORM_MUL_QUANTIZE=${FUSED_ADD_NORM_MUL_QUANTIZE:-1}
+export FUSED_SILU_W13=${FUSED_SILU_W13:-1}
+export FUSED_PAD_GRAD_ACCUM=${FUSED_PAD_GRAD_ACCUM:-1}
 export OFFLOAD_OPTIM=${OFFLOAD_OPTIM:-1}

 export DEFAULT_FLOAT="bfloat16" OPTIM_DTYPE="bfloat16"
@@ -35,7 +42,7 @@ export DATA_SEED=${DATA_SEED:-5760}
 export JITBEAM=${JITBEAM:-3}
 export BEAM_UOPS_MAX=6000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=1

-export FAKEDATA=1 BENCHMARK=10
+export FAKEDATA=${FAKEDATA:-1} BENCHMARK=${BENCHMARK:-10}
 if [ -z "$FULL_LAYERS" ]; then
  export LLAMA_LAYERS=2
 fi
@@ -2,7 +2,6 @@

 export PYTHONPATH="."
 export DEV=${DEV:-AMD}
-export EMULATE="AMD_CDNA4"
 export CHECK_OOB=0
 export REWRITE_STACK_LIMIT=5000000 HCQDEV_WAIT_TIMEOUT_MS=240000
 export DEVICE_IN_FUNCTION_BUG=1
@@ -10,12 +9,21 @@ export DEVICE_IN_FUNCTION_BUG=1
 export DEBUG=${DEBUG:-0}
 export HK_FLASH_ATTENTION=${HK_FLASH_ATTENTION:-1}
 export ALL2ALL=${ALL2ALL:-1}
+export LATE_ALLREDUCE=${LATE_ALLREDUCE:-0}
 export USE_ATOMICS=${USE_ATOMICS:-1}
 export ASM_GEMM=${ASM_GEMM:-1}
-export WQKV=${WQKV:-0}
+export WQKV=${WQKV:-1}
+export MASTER_WEIGHTS=${MASTER_WEIGHTS:-1}
+export FP8=${FP8:-1}
+export ALLREDUCE_CAST=${ALLREDUCE_CAST:-1}
+export FAST_CE=${FAST_CE:-1}
+export FUSED_INPUT_QUANTIZE=${FUSED_INPUT_QUANTIZE:-1}
+export FUSED_ADD_NORM_MUL_QUANTIZE=${FUSED_ADD_NORM_MUL_QUANTIZE:-1}
+export FUSED_SILU_W13=${FUSED_SILU_W13:-1}
+export FUSED_PAD_GRAD_ACCUM=${FUSED_PAD_GRAD_ACCUM:-1}

 export DEFAULT_FLOAT="bfloat16" OPTIM_DTYPE="bfloat16"
-export DP=${DP:-8} MP=${MP:-1} BS=${BS:-8} EVAL_BS=${EVAL_BS:-8} GRADIENT_ACC_STEPS=${GRADIENT_ACC_STEPS:-4}
+export DP=${DP:-8} MP=${MP:-1} BS=${BS:-16} EVAL_BS=${EVAL_BS:-8} GRADIENT_ACC_STEPS=${GRADIENT_ACC_STEPS:-2}
 export GBS=$((BS * GRADIENT_ACC_STEPS))

 export MODEL="llama3"
@@ -1,5 +1,6 @@
 #!/bin/bash
 export BENCHMARK=5
 export EVAL_BS=0
-VIZ=${VIZ:--1} examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama8b/implementations/tinybox_8xMI350X/dev_run.sh
-extra/viz/cli.py --profile --device "AMD" | head -23
+VIZ=${VIZ:--1} FULL_LAYERS=1 DEBUG=0 examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama8b/implementations/tinybox_8xMI350X/dev_beam.sh
+SRC="AMD"; [[ $DEV == NULL* ]] && SRC="NULL"
+python -m tinygrad.viz.cli -s "$SRC" -t
@@ -0,0 +1,55 @@
+#!/usr/bin/env bash
+set -e  # Exit on any error
+set -o pipefail  # Make pipeline fail if any command fails
+
+export PYTHONPATH="."
+export DEV=AMD
+export CHECK_OOB=0
+export REWRITE_STACK_LIMIT=5000000 HCQDEV_WAIT_TIMEOUT_MS=240000
+export DEVICE_IN_FUNCTION_BUG=1
+
+export HK_FLASH_ATTENTION=1
+export ALL2ALL=1
+export LATE_ALLREDUCE=0
+export USE_ATOMICS=1
+export ASM_GEMM=1
+export WQKV=1
+export MASTER_WEIGHTS=1
+export FP8=1
+export ALLREDUCE_CAST=1
+export FAST_CE=1
+export FUSED_INPUT_QUANTIZE=1
+export FUSED_ADD_NORM_MUL_QUANTIZE=1
+export FUSED_SILU_W13=1
+export FUSED_PAD_GRAD_ACCUM=1
+
+export DEFAULT_FLOAT="bfloat16" OPTIM_DTYPE="bfloat16"
+export DP=8 MP=1 BS=16 EVAL_BS=8 GRADIENT_ACC_STEPS=2
+export GBS=$((BS * GRADIENT_ACC_STEPS))
+
+export MODEL="llama3"
+export BASEDIR="/raid/datasets/c4-8b/"
+export SMALL=1
+export LLAMA3_SIZE=8B
+export EVAL_TARGET=3.3 EVAL_FREQ=12288
+export LR="1e-3" END_LR="1e-4" WARMUP_SAMPLES=4096 MAX_STEPS=1200000
+export WARMUP_STEPS=$((WARMUP_SAMPLES / GBS))
+export SAMPLES=$((MAX_STEPS * GBS))
+export SEQLEN=8192
+
+export SEED=$RANDOM
+export DATA_SEED=$SEED
+
+export JITBEAM=3
+export BEAM_UOPS_MAX=6000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=1
+
+export LOGMLPERF=1
+
+DATETIME=$(date "+%m%d%H%M")
+LOGFILE="llama31_8b_8xMI350x_${DATETIME}_${SEED}.log"
+
+# beam
+FAKEDATA=1 BENCHMARK=10 INITMLPERF=1 LLAMA_LAYERS=2 python3 examples/mlperf/model_train.py | tee "$LOGFILE"
+
+# run
+RUNMLPERF=1 python3 examples/mlperf/model_train.py | tee -a "$LOGFILE"
@@ -0,0 +1,38 @@
+{
+    "submitter": "tinycorp",
+    "division": "closed",
+    "status": "Available on-premise",
+    "system_name": "tinybox 8xMI350X",
+    "number_of_nodes": "1",
+    "host_processors_per_node": "2",
+    "host_processor_model_name": "AMD EPYC 9575F",
+    "host_processor_core_count": "32",
+    "host_processor_vcpu_count": "64",
+    "host_processor_frequency": "",
+    "host_processor_caches": "",
+    "host_processor_interconnect": "",
+    "host_memory_capacity": "3072 GiB",
+    "host_storage_type": "NVMe SSD",
+    "host_storage_capacity": "4TB",
+    "host_networking": "",
+    "host_networking_topology": "",
+    "host_memory_configuration": "24x 128GB DDR5",
+    "accelerators_per_node": "8",
+    "accelerator_model_name": "AMD Instinct MI350X 288GB HBM3e",
+    "accelerator_host_interconnect": "PCIe 5.0 x16",
+    "accelerator_frequency": "",
+    "accelerator_on-chip_memories": "",
+    "accelerator_memory_configuration": "HBM3",
+    "accelerator_memory_capacity": "288GB",
+    "accelerator_interconnect": "",
+    "accelerator_interconnect_topology": "",
+    "cooling": "air",
+    "hw_notes": "",
+    "framework": "tinygrad, branch mlperf_training_v6.0",
+    "other_software_stack": {
+        "python": "3.12.3",
+        "ROCm": "7.1.1"
+    },
+    "operating_system": "Ubuntu 24.04.3 LTS",
+    "sw_notes": ""
+  }
@@ -4,7 +4,7 @@ if "JIT_BATCH_SIZE" not in os.environ: os.environ["JIT_BATCH_SIZE"] = "0"

 from tinygrad import fetch, Tensor, TinyJit, Context, GlobalCounters, Device, dtypes
 from tinygrad.helpers import DEBUG, getenv
-from tinygrad.engine.realize import CompiledRunner
+from tinygrad.uop.ops import Ops
 from tinygrad.nn.onnx import OnnxRunner

 OPENPILOT_MODEL = sys.argv[1] if len(sys.argv) > 1 else "https://github.com/commaai/openpilot/raw/v0.9.7/selfdrive/modeld/models/supercombo.onnx"
@@ -35,7 +35,11 @@ def compile(onnx_file):
      ret = run_onnx_jit(**inputs).numpy()
    # copy i == 1 so use of JITBEAM is okay
    if i == 1: test_val = np.copy(ret)
-  print(f"captured {len(run_onnx_jit.captured.jit_cache)} kernels")
+  # iterate kernel CALLs in the captured LINEAR UOp; toposort descends into batched graph CUSTOM_FUNCTIONs
+  kernel_asts = {Ops.PROGRAM}
+  kernel_calls = [u for u in run_onnx_jit.captured.linear.toposort(gate=lambda x: x.op not in kernel_asts)
+                  if u.op is Ops.CALL and u.src[0].op in kernel_asts]
+  print(f"captured {len(kernel_calls)} kernels")
  np.testing.assert_equal(test_val, ret, "JIT run failed")
  print("jit run validated")

@@ -43,13 +47,14 @@ def compile(onnx_file):
  kernel_count = 0
  read_image_count = 0
  gated_read_image_count = 0
-  for ei in run_onnx_jit.captured.jit_cache:
-    if isinstance(ei.prg, CompiledRunner):
-      kernel_count += 1
-      read_image_count += ei.prg.p.src.count("read_image")
-      gated_read_image_count += ei.prg.p.src.count("?read_image")
-      for v in [m.group(1) for m in re.finditer(r'(val\d+)\s*=\s*read_imagef\(', ei.prg.p.src)]:
-        if len(re.findall(fr'[\?\:]{v}\.[xyzw]', ei.prg.p.src)) > 0: gated_read_image_count += 1
+  for call in kernel_calls:
+    _, _, _, source, _ = call.src[0].src
+    src = source.arg
+    kernel_count += 1
+    read_image_count += src.count("read_image")
+    gated_read_image_count += src.count("?read_image")
+    for v in [m.group(1) for m in re.finditer(r'(val\d+)\s*=\s*read_imagef\(', src)]:
+      if len(re.findall(fr'[\?\:]{v}\.[xyzw]', src)) > 0: gated_read_image_count += 1
  print(f"{kernel_count=},  {read_image_count=}, {gated_read_image_count=}")
  if (allowed_kernel_count:=getenv("ALLOWED_KERNEL_COUNT", -1)) != -1:
    assert kernel_count == allowed_kernel_count, f"different kernels! {kernel_count=}, {allowed_kernel_count=}"
@@ -128,14 +133,20 @@ def bench(run, inputs):
      run(**inputs).numpy()

 if __name__ == "__main__":
-  onnx_file = fetch(OPENPILOT_MODEL)
-  inputs, outputs = compile(onnx_file)
+  if getenv("RUN_PICKLE"):
+    with open(OUTPUT, "rb") as f: pickle_loaded = pickle.load(f)
+    inputs = {name: Tensor(Tensor.randn(*[int(s) for s in view.src[1].arg], dtype=dtype).numpy(), device=device)
+              for name, (view, _vars, dtype, device) in zip(pickle_loaded.captured.expected_names, pickle_loaded.captured.expected_input_info)}
+    test_vs_compile(pickle_loaded, inputs)
+  else:
+    onnx_file = fetch(OPENPILOT_MODEL)
+    inputs, outputs = compile(onnx_file)

-  with open(OUTPUT, "rb") as f: pickle_loaded = pickle.load(f)
+    with open(OUTPUT, "rb") as f: pickle_loaded = pickle.load(f)

-  test_vs_compile(pickle_loaded, inputs, outputs)
-  if getenv("SELFTEST"):
-    test_vs_onnx(inputs, outputs, onnx_file, 1e-4)
+    test_vs_compile(pickle_loaded, inputs, outputs)
+    if getenv("SELFTEST"):
+      test_vs_onnx(inputs, outputs, onnx_file, 1e-4)

  if getenv("BENCHMARK_LOG", ""):
    bench(pickle_loaded, inputs)
@@ -66,7 +66,7 @@ if __name__ == "__main__":
  model_path = Path(args.weights) if args.weights else download_weights(model_info["total_num_weights"])
  transformer = load_model(model_path, model_info["model_params"])
  tokenizer = AutoTokenizer.from_pretrained(model_info["tokenizer"])
-  param_bytes = sum(x.uop.size * x.dtype.itemsize for x in get_parameters(transformer))
+  param_bytes = sum(x.nbytes() for x in get_parameters(transformer))

  outputted = args.prompt
  start_pos, toks = 0, tokenizer(outputted)["input_ids"]
@@ -3,7 +3,7 @@ from extra.export_model import export_model
 from examples.llama3 import build_transformer, Tokenizer
 from tinygrad.nn.state import get_state_dict, load_state_dict
 from tinygrad import Device, Variable, Tensor, dtypes, TinyJit
-from tinygrad.helpers import fetch, Context
+from tinygrad.helpers import DEV, fetch, Context
 from tiktoken.load import load_tiktoken_bpe, dump_tiktoken_bpe

 def prepare_browser_chunks(model):
@@ -115,7 +115,7 @@ if __name__=="__main__":
  start_pos = Variable("start_pos", 0, max_context).bind(0)
  model_input = lambda: [Tensor([[tok]]), start_pos, TEMPERATURE, TOP_K, TOP_P, ALPHA_F, ALPHA_P]

-  Device.DEFAULT="CPU"
+  DEV.value = "CPU"
  model = build_transformer(model_path, model_size="1B", quantize="int8", scale_dtype=dtypes.float32, device=Device.DEFAULT, max_context=max_context)
  state_dict = get_state_dict(model)
  validate_model(model, tokenizer)
@@ -129,7 +129,7 @@ if __name__=="__main__":
  with open(os.path.join(os.path.dirname(__file__), f"{model_name}.c"), "w") as f: f.write(cprog)
  with open(os.path.join(os.path.dirname(__file__), "net_clang.js"), "w") as f: f.write(js_wrapper)

-  Device.DEFAULT="WEBGPU"
+  DEV.value = "WEBGPU"
  # float16 is not yet supported for dawn/Vulkan/NVIDIA stack, see: https://issues.chromium.org/issues/42251215
  # therefore for now, we used CLANG to quantize the float16 llama to int8 with float32 scales, then load to WEBGPU
  model = build_transformer(model_path, model_size="1B", quantize="int8", max_context=max_context, load_weights=False)
@@ -1,7 +1,6 @@
 from tinygrad import Tensor, Device, TinyJit, dtypes
-from tinygrad.helpers import getenv

-GPUS = getenv("GPUS", 4) # TODO: expose a way in tinygrad to access this
+GPUS = Device[Device.DEFAULT].count()
 N = 6144

@TinyJit
@@ -4,8 +4,8 @@ from extra.f16_decompress import u32_to_f16
 from examples.stable_diffusion import StableDiffusion
 from tinygrad.nn.state import get_state_dict, safe_save, safe_load_metadata, torch_load, load_state_dict
 from tinygrad.tensor import Tensor
-from tinygrad import Device, dtypes
-from tinygrad.helpers import fetch
+from tinygrad import dtypes
+from tinygrad.helpers import DEV, fetch
 from typing import NamedTuple, Any, List
 import requests
 import argparse
@@ -80,7 +80,7 @@ if __name__ == "__main__":
  parser = argparse.ArgumentParser(description='Run Stable Diffusion', formatter_class=argparse.ArgumentDefaultsHelpFormatter)
  parser.add_argument('--remoteweights', action='store_true', help="Use safetensors from Huggingface, or from local")
  args = parser.parse_args()
-  Device.DEFAULT = "WEBGPU"
+  DEV.value = "WEBGPU"

  model = StableDiffusion()

@@ -111,19 +111,19 @@ if __name__ == "__main__":
    return code

  def compile_step(model, step: Step):
-    run, special_names = jit_model(step, *step.input)
-    functions, statements, bufs, _ = compile_net(run, special_names)
+    linear, output_bufs = jit_model(step, *step.input)
+    functions, statements, bufs, _ = compile_net(linear, output_bufs)
    state = get_state_dict(model)
-    weights = {id(x.uop.base.realized): name for name, x in state.items()}
+    weights = {(id(b), b.offset, b.size, b.dtype): name for name, x in state.items() if (b:=x.uop.base.realized) is not None}
    kernel_code = '\n\n'.join([f"const {key} = `{fixup_code(code, key)}`;" for key, code in functions.items()])
    kernel_names = ', '.join([name for (name, _, _, _) in statements])
-    input_names = [name for _,name in special_names.items() if "input" in name]
-    output_names = [name for _,name in special_names.items() if "output" in name]
+    input_names = [f"input{i}" for i in range(len(step.input))]
+    output_names = [f"output{i}" for i in range(len(output_bufs))]
    input_buf_types = [dtype_to_js_type(bufs[inp_name][1]) for inp_name in input_names]
    output_buf_types = [dtype_to_js_type(bufs[out_name][1]) for out_name in output_names]
    kernel_calls = '\n        '.join([f"addComputePass(device, commandEncoder, piplines[{i}], [{', '.join(args)}], {global_size});" for i, (_name, args, global_size, _local_size) in enumerate(statements) ])
    exported_bufs =  '\n    '.join([f"const {name} = " + (f"createEmptyBuf(device, {size});" if _key not in weights else f"createWeightBuf(device, {size}, getTensorBuffer(safetensor, metadata['{weights[_key]}'], '{weights[_key]}'))") + ";"  for name,(size,dtype,_key) in bufs.items()])
-    gpu_write_bufs =  '\n    '.join([f"const gpuWriteBuffer{i} = device.createBuffer({{size:input{i}.size, usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.MAP_WRITE }});" for i,(_,value) in enumerate(special_names.items()) if "output" not in value])
+    gpu_write_bufs =  '\n    '.join([f"const gpuWriteBuffer{i} = device.createBuffer({{size:input{i}.size, usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.MAP_WRITE }});" for i in range(len(input_names))])
    input_writer = '\n    '.join([f"await gpuWriteBuffer{i}.mapAsync(GPUMapMode.WRITE);\n    new {input_buf_types[i]}(gpuWriteBuffer{i}.getMappedRange()).set(" + f'data{i});' + f"\n    gpuWriteBuffer{i}.unmap();\ncommandEncoder.copyBufferToBuffer(gpuWriteBuffer{i}, 0, input{i}, 0, gpuWriteBuffer{i}.size);"  for i,_ in enumerate(input_names)])
    return f"""\n    var {step.name} = function() {{

@@ -141,7 +141,7 @@ if __name__ == "__main__":
        const kernels = [{kernel_names}];
        const piplines = await Promise.all(kernels.map(name => device.createComputePipelineAsync({{layout: "auto", compute: {{ module: device.createShaderModule({{ code: name }}), entryPoint: "main" }}}})));

-        return async ({",".join([f'data{i}' for i,(k,v) in enumerate(special_names.items()) if v != "output0"])}) => {{
+        return async ({",".join([f'data{i}' for i in range(len(input_names))])}) => {{
            const commandEncoder = device.createCommandEncoder();

            {input_writer}
@@ -4,10 +4,11 @@ from tinygrad.tensor import Tensor
 from tinygrad.nn.state import safe_save
 from extra.export_model import export_model
 from tinygrad.device import Device
+from tinygrad.helpers import DEV
 from tinygrad.nn.state import safe_load, load_state_dict

 if __name__ == "__main__":
-    Device.DEFAULT = "WEBGPU"
+    DEV.value = "WEBGPU"
    yolo_variant = 'n'
    yolo_infer = YOLOv8(w=0.25, r=2.0, d=0.33, num_classes=80)
    state_dict = safe_load(get_weights_location(yolo_variant))
@@ -64,7 +64,7 @@ def get_bar0_size(pcibus):

 class AMSMI(AMDev):
  def __init__(self, pcibus, vram_bar:MMIOInterface, doorbell_bar:MMIOInterface, mmio_bar:MMIOInterface):
-    self.pcibus = pcibus
+    self.pcibus, self.devfmt = pcibus, pcibus
    self.vram, self.doorbell64, self.mmio = vram_bar, doorbell_bar, mmio_bar
    self.pci_state = self.read_pci_state()
    if self.pci_state == "D0": self._init_from_d0()
@@ -91,6 +91,7 @@ class SMICtx:
    self.prev_lines_cnt = 0
    self.prev_terminal_width = 0
    self.prev_terminal_height = 0
+    self.prev_metrics = {}

    remove_parts = ["Advanced Micro Devices, Inc. [AMD/ATI]", "VGA compatible controller:", "Processing accelerators:"]
    lspci = subprocess.check_output(["lspci"]).decode("utf-8").splitlines()
@@ -235,6 +236,29 @@ class SMICtx:
      case (13,0,12): return self._smuq10_round(metrics.SocketPower), self._smuq10_round(metrics.SocketPowerLimit)
      case _: return metrics.SmuMetrics.AverageSocketPower, metrics.SmuMetrics.dGPU_W_MAX

+  def get_throttle_info(self, dev, metrics):
+    match dev.ip_ver[am.MP1_HWIP]:
+      case (13,0,6)|(13,0,12):
+        throttle_fields = [('ProchotResidencyAcc', 'Prochot'), ('PptResidencyAcc', 'PPT'),
+                           ('SocketThmResidencyAcc', 'Socket Thm'), ('VrThmResidencyAcc', 'VR Thm'), ('HbmThmResidencyAcc', 'HBM Thm')]
+        prev = self.prev_metrics.get(dev.pcibus)
+        active = []
+        if prev is not None:
+          acc_delta = metrics.AccumulationCounter - prev.AccumulationCounter
+          if acc_delta > 0:
+            for field, name in throttle_fields:
+              delta = getattr(metrics, field) - getattr(prev, field)
+              if delta > 0 and (pct := min(100, (delta * 100 + acc_delta // 2) // acc_delta)) > 0: active.append((name, pct))
+        return active
+      case _:
+        smu_mod = dev.smu.smu_mod
+        throttler_names = {getattr(smu_mod, a): a[len('THROTTLER_'):-len('_BIT')]
+                           for a in dir(smu_mod) if a.startswith('THROTTLER_') and a.endswith('_BIT')}
+        active = []
+        for i, pct in enumerate(metrics.SmuMetrics.ThrottlingPercentage):
+          if pct > 0: active.append((throttler_names.get(i, f"UNK_{i}"), int(pct)))
+        return active
+
  def get_mem_usage(self, dev):
    usage = 0
    pt_stack = [dev.mm.root_page_table]
@@ -281,6 +305,13 @@ class SMICtx:
                    + [f"MEM Activity {draw_bar(self.get_mem_activity(dev, metrics) / 100, activity_line_width)}"] \
                    + [f"MEM Usage    {draw_bar(mem_used / mem_total, activity_line_width, opt_text=mem_fmt)}"] \

+      throttle_info = self.get_throttle_info(dev, metrics)
+      if throttle_info:
+        throttle_text = colored(', '.join(f"{name} {pct}%" for name, pct in throttle_info), "red")
+      else:
+        throttle_text = colored("None", "green")
+      activity_line += [f"Throttle     {throttle_text}" + " " * (activity_line_width + 2)]
+
      temps_data, temps_data_compact = self.get_temps(dev, metrics), self.get_temps(dev, metrics, compact=True)
      temps_table = ["=== Temps (°C) ==="] + [f"{name:<16}: {color_temp(val)}" for name, val in temps_data.items()]
      temps_table_compact = ["Temps (°C):" + '/'.join([f"{color_temp(val)} {name}" for name, val in temps_data_compact.items()])]
@@ -324,6 +355,8 @@ class SMICtx:

      dev_content.append(device_line + activity_line + same_line([temps_table, power_table, frequency_table]))

+    self.prev_metrics = {dev.pcibus: m for dev, m in dev_metrics.items() if m is not None}
+
    raw_text = 'AM Monitor'.center(terminal_width) + "\n" + "=" * terminal_width + "\n\n"
    for i in range(0, len(dev_content), 2):
      if i + 1 < len(dev_content): raw_text += '\n'.join(same_line([dev_content[i], dev_content[i+1]], split=padding))
@@ -28,15 +28,7 @@
 // #include "soc15_ih_clientid.h"
 // #include "amdgpu_ih.h"

-#define int32_t int
-#define uint32_t unsigned int
-#define int8_t signed char
-#define uint8_t unsigned char
-#define uint16_t unsigned short
-#define int16_t short
-#define uint64_t unsigned long long
 #define bool _Bool
-#define u32 unsigned int

 #define AMDGPU_MAX_IRQ_SRC_ID		0x100
 #define AMDGPU_MAX_IRQ_CLIENT_ID	0x100
@@ -22,15 +22,7 @@
 #ifndef __AMDGPU_SMU_H__
 #define __AMDGPU_SMU_H__

-#define int32_t int
-#define uint32_t unsigned int
-#define int8_t signed char
-#define uint8_t unsigned char
-#define uint16_t unsigned short
-#define int16_t short
-#define uint64_t unsigned long long
 #define bool _Bool
-#define u32 unsigned int

 #define SMU_THERMAL_MINIMUM_ALERT_TEMP		0
 #define SMU_THERMAL_MAXIMUM_ALERT_TEMP		255
@@ -24,15 +24,7 @@
 #define __AMDGPU_UCODE_H__

 // #include "amdgpu_socbb.h"
-#define int32_t int
-#define uint32_t unsigned int
-#define int8_t signed char
-#define uint8_t unsigned char
-#define uint16_t unsigned short
-#define int16_t short
-#define uint64_t unsigned long long
 #define bool _Bool
-#define u32 unsigned int

 struct common_firmware_header {
 	uint32_t size_bytes; /* size of the entire header+image(s) in bytes */
@@ -1,47 +1,50 @@
 from typing import Tuple, Dict, List, Optional
 from tinygrad.dtype import DType, dtypes
-from tinygrad.renderer import ProgramSpec
 from tinygrad.tensor import Tensor
-from tinygrad.device import Device
+from tinygrad.device import Device, Buffer
 from tinygrad.engine.jit import TinyJit
 from tinygrad.nn.state import get_state_dict
-from tinygrad.helpers import Context, to_mv
-from tinygrad.uop.ops import Ops
+from tinygrad.helpers import Context, to_mv, prod
+from tinygrad.uop.ops import Ops, UOp
+from tinygrad.codegen import to_program
 import json
 from collections import OrderedDict

 EXPORT_SUPPORTED_DEVICE = ["WEBGPU", "CPU", "CUDA", "CL"]

-def compile_net(run:TinyJit, special_names:Dict[int,str]) -> Tuple[Dict[str,str],List[Tuple[str,List[str],List[int]]],Dict[str,Tuple[int,DType,int]],Dict[str,Tensor]]:
-  # memory-planned subbuffers can have multiple Buffer objects for the same memory region
-  canon, _seen = {}, {}
-  for ji in run.jit_cache:
-    for b in ji.bufs:
-      if b is not None: canon[id(b)] = _seen.setdefault((id(b.base._buf), b.offset, b.size, b.dtype), b)
-  special_names = {id(canon[k]): v for k, v in special_names.items() if k in canon}
+_KERNEL_ASTS = {Ops.SINK, Ops.PROGRAM}
+def iter_kernel_calls(linear:UOp):
+  """Yield kernel CALLs from a LINEAR UOp. Toposort descends naturally into CUSTOM_FUNCTION graph batches; gate stops at kernel ASTs."""
+  return (u for u in linear.toposort(gate=lambda x: x.op not in _KERNEL_ASTS) if u.op is Ops.CALL and u.src[0].op in _KERNEL_ASTS)

-  functions, bufs, bufs_to_save, statements, bufnum = {}, {}, {}, [], 0
-  for ji in run.jit_cache:
-    fxn: ProgramSpec = ji.prg.p
-    functions[fxn.function_name] = fxn.src   # NOTE: this assumes all with the same name are the same
-    cargs = []
-    for i,arg in enumerate(ji.bufs):
-      arg = canon[id(arg)]
-      key = id(arg)
-      if key not in bufs:
-        if key in special_names:
-          bufs[key] = (special_names[key], arg.size*arg.dtype.itemsize, arg.dtype, key)
-        else:
-          bufs[key] = (f"buf_{bufnum}", arg.size*arg.dtype.itemsize, arg.dtype, key)
-          bufnum += 1
-          if i > 0: bufs_to_save[bufs[key][0]] = arg   # if first usage of a buffer is not an output, and it's not a special name
-      cargs.append(bufs[key][0])
-    cargs += [var for var in fxn.vars if getattr(var, "op", None) is Ops.DEFINE_VAR] # symbolic vars; is it necessary or sufficient to check for DEFINE_VAR?
-    statements.append((fxn.function_name, cargs, fxn.global_size, fxn.local_size))
+def compile_net(linear:UOp, output_bufs:List[Buffer]) -> Tuple[Dict[str,str], List, Dict[str,Tuple[int,DType,int]], Dict[str,Buffer]]:
+  output_name = {id(b): f"output{i}" for i, b in enumerate(output_bufs)}
+  functions, bufs, bufs_to_save, statements, n = {}, {}, {}, [], 0

-  return functions, statements, {name:(size, dtype, key) for (name,size,dtype,key) in bufs.values()}, bufs_to_save
+  def name_of(bu:UOp, is_out:bool) -> str:
+    nonlocal n
+    if bu.op is Ops.PARAM: key, name, size = ("in", bu.arg), f"input{bu.arg}", prod(bu.shape)*bu.dtype.itemsize
+    else:
+      b = bu.buffer
+      key, size = (id(b.base), b.offset, b.size, b.dtype), b.size*b.dtype.itemsize
+      if key in bufs: return bufs[key][0]
+      if (name:=output_name.get(id(b))) is None:
+        name, n = f"buf_{n}", n+1
+        if not is_out: bufs_to_save[name] = b
+    bufs[key] = (name, size, bu.dtype, key)
+    return name

-def jit_model(model, *args) -> Tuple[TinyJit,Dict[int,str]]:
+  for call in iter_kernel_calls(linear):
+    arg_uops = [b for b in call.src[1:] if b.op is not Ops.BIND]
+    prg = to_program(call.src[0], Device[arg_uops[0].device].renderer)
+    info = prg.arg
+    functions[info.function_name] = prg.src[3].arg
+    cargs = [name_of(bu, i == 0) for i, bu in enumerate(arg_uops)] + [v for v in info.vars if v.op is Ops.DEFINE_VAR]
+    statements.append((info.function_name, cargs, info.global_size, info.local_size))
+
+  return functions, statements, {name:(size, dtype, key) for name, size, dtype, key in bufs.values()}, bufs_to_save
+
+def jit_model(model, *args) -> Tuple[UOp, List[Buffer]]:
  assert hasattr(model, "forward") or callable(model), "model needs a forward function"
  @TinyJit
  def run(*x):
@@ -50,20 +53,10 @@ def jit_model(model, *args) -> Tuple[TinyJit,Dict[int,str]]:
    out = [out] if isinstance(out, Tensor) else out
    return [o.realize() for o in out]

-  # twice to run the JIT
+  # run twice to trigger JIT capture
  for _ in range(2): the_output = run(*args)
-  special_names = {}
-
-  # hack to put the inputs back
-  for (j,i),idx in run.input_replace.items():
-    realized_input = args[idx].uop.base.realized
-    run.jit_cache[j].bufs[i] = realized_input
-    special_names[id(realized_input)] = f'input{idx}'
-
-  # TODO: fetch this from the jit in self.input_replace and self.ret (hint: use get_parameters on self.ret)
-  for i, output in enumerate(the_output):
-    special_names[id(output.uop.base.realized)] = f'output{i}'
-  return run, special_names
+  assert run.captured is not None
+  return run.captured.linear, [o.uop.base.realized for o in the_output]

 def export_model_clang(functions:Dict[str,str], statements:Dict[str,Tuple[str,int,int]], bufs:Dict[str,Tuple[str,int,int]],
  bufs_to_save:Dict[str,Tensor], input_names:List[str], output_names:List[str], weight_names={}, model_name="model", symbolic_vars={}, wasm=False) -> str:
@@ -249,12 +242,12 @@ def export_model(model, target:str, *inputs, model_name: Optional[str] = "model"
  assert Device.DEFAULT in EXPORT_SUPPORTED_DEVICE, f"only {', '.join(EXPORT_SUPPORTED_DEVICE)} are supported"

  # NOTE: CPU_COUNT=1, since export does not support threading
-  with Context(JIT=2, CPU_COUNT=1): run,special_names = jit_model(model, *inputs)
-  functions, statements, bufs, bufs_to_save = compile_net(run, special_names)
+  with Context(JIT=2, CPU_COUNT=1): linear, output_bufs = jit_model(model, *inputs)
+  functions, statements, bufs, bufs_to_save = compile_net(linear, output_bufs)
  state = get_state_dict(model)
-  weight_names = {id(x.uop.base.realized): name for name, x in state.items()}
-  input_names = [name for _,name in special_names.items() if "input" in name]
-  output_names = [name for _,name in special_names.items() if "output" in name]
+  weight_names = {(id(b), b.offset, b.size, b.dtype): name for name, x in state.items() if (b:=x.uop.base.realized) is not None}
+  input_names = [f"input{i}" for i in range(len(inputs))]
+  output_names = [f"output{i}" for i in range(len(output_bufs))]

  # handle symbolic variables; TODO: refactor to fix some of this stuff upstream in tinygrad
  symbolic_vars = OrderedDict()
@@ -13,7 +13,7 @@ from tinygrad import Tensor, Device, Context, GlobalCounters
 from tinygrad.uop.ops import UOp, Ops, KernelInfo
 from tinygrad.helpers import getenv, colored
 from tinygrad.dtype import dtypes, AddrSpace
-from tinygrad.engine.realize import Estimates
+from tinygrad.engine.realize import Estimates, run_linear
 from tinygrad.renderer.amd.dsl import s, v, VCC_LO, NULL
 from tinygrad.runtime.autogen.amd.rdna3.ins import *

@@ -167,7 +167,7 @@ PREFETCH_LOADS = [(V_LDS_A_DATA[4+2*i], V_LDS_A_DATA[4+2*i+1], V_GLOBAL_B_ADDR,
 # =============================================================================

 class Kernel:
-  def __init__(self, arch='gfx1100'): self.instructions, self.labels, self.pos, self.arch = [], {}, 0, arch
+  def __init__(self): self.instructions, self.labels, self.pos = [], {}, 0
  def label(self, name): self.labels[name] = self.pos

  def emit(self, inst, target=None):
@@ -196,10 +196,10 @@ class Kernel:
 # Kernel builder
 # =============================================================================

-def build_kernel(N, arch='gfx1100'):
+def build_kernel(N):
  assert N % 128 == 0, f"N must be a multiple of 128 (tile size), got {N}"
  assert N >= 256, f"N must be >= 256 (prefetch pipeline requires at least 2 K-blocks), got {N}"
-  k = Kernel(arch)
+  k = Kernel()

  # ===========================================================================
  # PROLOGUE: Load kernel arguments, compute tile coordinates and addresses
@@ -441,9 +441,9 @@ THREADS = 128

 def test_matmul():
  dev = Device[Device.DEFAULT]
-  print(f"Device arch: {dev.renderer.arch}")
+  print(f"Device arch: {dev.renderer.target.arch}")

-  insts = build_kernel(N, dev.renderer.arch)
+  insts = build_kernel(N)

  rng = np.random.default_rng(42)
  a = Tensor(rng.random((N, N), dtype=np.float32) - 0.5)
@@ -463,11 +463,14 @@ def test_matmul():
                                                                                  estimates=Estimates(ops=N*N*N*2, mem=N*N*4*3)))
    return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=tuple([UOp(Ops.INS, arg=x) for x in insts]))))
  c = Tensor.custom_kernel(a, b, c, fxn=asm_kernel)[2]
-  ei = c.schedule()[0].lower()
+  linear = c.schedule_linear()

  ets = []
  with Context(DEBUG=2):
-    for _ in range(getenv("CNT", 5)): ets.append(ei.run(wait=True))
+    for _ in range(getenv("CNT", 5)):
+      start = GlobalCounters.time_sum_s
+      run_linear(linear)
+      ets.append(GlobalCounters.time_sum_s - start)
  print(f"REAL TFLOPS {N * N * N * 2 / min(ets) * 1e-12:.2f}")

  if getenv("VERIFY", 1):
@@ -1,4 +1,4 @@
-from tinygrad import UOp, getenv
+from tinygrad import Device, UOp, getenv
 from tinygrad.uop.ops import AxisType, KernelInfo, Ops
 from tinygrad.dtype import AddrSpace, dtypes

@@ -13,18 +13,23 @@ assert N % BLOCK_N == 0 and M % BLOCK_M == 0 and K % BLOCK_K == 0

 use_wmma = getenv("WMMA")
 if use_wmma:
+  is_rdna4 = Device[Device.DEFAULT].renderer.target.arch.startswith("gfx12")
+
  WAVES_M, WAVES_N = 2, 2
  LANES_PER_WAVE_M, LANES_PER_WAVE_N = 2, 16
-  UNROLL_M, UNROLL_N = 1, 1

  # wmma params
  WMMA_M, WMMA_N, WMMA_K = 16, 16, 16
  WMMA_ACC = WMMA_M // LANES_PER_WAVE_M
+  UNROLL_M, UNROLL_N = (WMMA_ACC, 1) if is_rdna4 else (1, 1)
 else:
  WAVES_M, WAVES_N = 4, 1
  LANES_PER_WAVE_M, LANES_PER_WAVE_N = 4, 8
  UNROLL_M, UNROLL_N = 4, 4

+# total lanes must be the warp size
+assert LANES_PER_WAVE_M*LANES_PER_WAVE_N == WARP_SIZE
+
 # WARP_SIZE * total waves
 THREADS_PER_BLOCK = WARP_SIZE * WAVES_M * WAVES_N

@@ -61,7 +66,7 @@ def block_128x128_gemm(c:UOp, a:UOp, b:UOp) -> UOp:

  # accumulator (unified: both paths use (TM, TN) with scalar dtypes.float)
  acc = UOp.placeholder((TM, TN), dtypes.float, slot=2, addrspace=AddrSpace.REG)
-  acc = acc.after(acc.store(UOp.const(dtypes.float, 0).reshape((1,)*len(acc.shape)).expand(acc.shape)))
+  acc = acc.after(acc.store(acc.zeros_like()))

  if use_wmma:
    k = UOp.range(BLOCK_K // WMMA_K, 101, AxisType.REDUCE)
@@ -71,7 +76,10 @@ def block_128x128_gemm(c:UOp, a:UOp, b:UOp) -> UOp:
    acc_frag = acc.reshape(TM // WMMA_ACC, WMMA_ACC, TN).permute(0,2,1)[tile_m, tile_n]
    a_frag = A_local.reshape(WAVES_M, TM // WMMA_ACC, WMMA_M, BLOCK_K // WMMA_K, WMMA_K)[wave_m, tile_m, lane_n, k]
    b_frag = B_local.reshape(WAVES_N, TN, WMMA_N, BLOCK_K // WMMA_K, WMMA_K)[wave_n, tile_n, lane_n, k]
-
+    if is_rdna4:
+      # NOTE: since this is part of K, these 2 can be anywhere in the frags and long as a and b match
+      a_frag = a_frag.reshape(2, 8)[lane_m, :]
+      b_frag = b_frag.reshape(2, 8)[lane_m, :]
    wmma = UOp(Ops.SHAPED_WMMA, dtypes.float, (a_frag, b_frag, acc_frag.after(k)), arg=((16, 16, 16), 'AMD', 32))
    acc_store = acc_frag.store(wmma).end(tile_m, tile_n)
  else:
@@ -1,31 +1,39 @@
 # kernel8_batched_gmem.s from https://seb-v.github.io/optimization/update/2025/01/20/Fast-GPU-Matrix-multiplication.html
 # sudo PATH=/opt/homebrew/Cellar/llvm/20.1.6/bin:$PATH AMD_LLVM=0 AMD=1 DEBUG=2 python3 extra/gemm/amd_matmul.py
 import pathlib
-from dataclasses import replace
 from tinygrad import Tensor, Device, Context, GlobalCounters
 from tinygrad.helpers import getenv
-from tinygrad.engine.realize import CompiledRunner, ExecItem, get_program
+from tinygrad.uop.ops import UOp, Ops, KernelInfo
+from tinygrad.renderer import Estimates
+from tinygrad.engine.realize import run_linear

 N = 4096
 run_count = 5

-if __name__ == "__main__":
-  ast = (Tensor.empty(N, N)@Tensor.empty(N, N)).schedule()[-1].ast
-  prg = get_program(ast, Device.default.renderer)
+def make_matmul_kernel(name:str, src:str, local_size:int):
+  def fxn(a:UOp, b:UOp, c:UOp) -> UOp:
+    threads = UOp.special(local_size, "lidx0")
+    wg_x = UOp.special(N//128, "gidx0")
+    wg_y = UOp.special(N//128, "gidx1")
+    sink = UOp.sink(a.base, b.base, c.base, threads, wg_x, wg_y, arg=KernelInfo(name, estimates=Estimates(ops=2*N**3, mem=3*N*N*4)))
+    lib = Device[Device.DEFAULT].compiler.compile_cached(src)
+    return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=Device.DEFAULT), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                                 UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=lib)))
+  return fxn

+if __name__ == "__main__":
  if getenv("ASM") == 1:
    src = (pathlib.Path(__file__).parent / "amd_seb" / "kernel8_batched_gmem.s").read_text()
-    prgfast = replace(prg, name="kernel", src=src, global_size=[N//128, N//128, 1], local_size=[128, 1, 1])
+    name, local_size = "kernel", 128
  elif getenv("ASM") == -1:
    src = (pathlib.Path(__file__).parent / "amd_seb" / "kernel3_registers.cpp").read_text()
-    prgfast = replace(prg, name="kernel3_registers", src=src, global_size=[N//128, N//128, 1], local_size=[256, 1, 1])
+    name, local_size = "kernel3_registers", 256
  elif getenv("ASM") == -2:
    src = (pathlib.Path(__file__).parent / "amd_seb" / "kernel4_gmem_df.cpp").read_text()
-    prgfast = replace(prg, name="kernel4_gmem_db", src=src, global_size=[N//128, N//128, 1], local_size=[256, 1, 1])
+    name, local_size = "kernel4_gmem_db", 256
  else:
    src = (pathlib.Path(__file__).parent / "amd_seb" / "kernel5_lds_optim.cpp").read_text()
-    prgfast = replace(prg, name="kernel5_lds_optim", src=src, global_size=[N//128, N//128, 1], local_size=[128, 1, 1])
-  runner = CompiledRunner(prgfast)
+    name, local_size = "kernel5_lds_optim", 128

  a = Tensor.randn(N, N).realize()
  b = Tensor.randn(N, N).realize()
@@ -35,8 +43,8 @@ if __name__ == "__main__":
  with Context(DEBUG=2):
    for _ in range(run_count): tc = (a@b).realize()

+  linear = Tensor.custom_kernel(a, b, c, fxn=make_matmul_kernel(name, src, local_size))[2].schedule_linear()
  GlobalCounters.reset()
-  ei = ExecItem(ast, [a.uop.buffer, b.uop.buffer, c.uop.buffer], prg=runner)
  with Context(DEBUG=2):
-    for _ in range(run_count): ei.run(wait=True)
+    for _ in range(run_count): run_linear(linear)
  print(f"custom  {(c-tc).square().mean().item()}")
@@ -44,9 +44,9 @@ nc = np.random.randn(N, N).astype(np.float32)

 ns = nb.reshape(-1, 32).sum(axis=0)

-a = MallocAllocator.alloc(na.size * np.dtype(np.float32).itemsize)
-b = MallocAllocator.alloc(nb.size * np.dtype(np.float32).itemsize)
-c = MallocAllocator.alloc(nc.size * np.dtype(np.float32).itemsize)
+a = MallocAllocator.alloc(na.nbytes)
+b = MallocAllocator.alloc(nb.nbytes)
+c = MallocAllocator.alloc(nc.nbytes)

 MallocAllocator._copyin(b, flat_mv(nb.data))
 MallocAllocator._copyin(c, flat_mv(nc.data))
@@ -1,10 +1,12 @@
-import atexit, functools
+import atexit, functools, pathlib
 from tinygrad import Tensor, Device, dtypes
 from tinygrad.dtype import AddrSpace
 from tinygrad.uop.ops import UOp, Ops, KernelInfo, AxisType
 from tinygrad.renderer import Estimates
 from tinygrad.helpers import getenv, all_same, DEBUG
+from tinygrad.runtime.support.compiler_amd import HIPCCCompiler
 from tinygrad.runtime.autogen.amd.cdna.ins import *
+from examples.mlperf.models.flat_llama import FP8_DTYPE, FP8_GRAD_DTYPE, quantize_fp8

 # ** CDNA4 assembly gemm

@@ -2623,6 +2625,30 @@ def custom_asm_gemm(C:UOp, A:UOp, B:UOp, dname:str) -> UOp:
  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname),
                                UOp(Ops.LINEAR, src=tuple([UOp(Ops.INS, arg=x) for x in insts]))))

+# ** FP8 GEMM custom kernel
+
+@functools.cache
+def custom_hk_fp8_gemm(C:UOp, A:UOp, B:UOp, *args:UOp, dname:str, scale_mode:int=3) -> UOp:
+  # scale_mode: 0=no scale, 1=x only, 2=w only, 3=both
+  n_scales = (1 if scale_mode & 1 else 0) + (1 if scale_mode & 2 else 0)
+  scales, extra = args[:n_scales], args[n_scales:]
+  M, K = A.shape[0]*A.shape[1], A.shape[2]
+  N, K2 = B.shape[(1 if B.ndim == 3 else 0):]
+  assert K == K2, f"{A.shape} {B.shape}"
+  block_size = 256
+  threads = UOp.special(64 * 8, "lidx0")
+  workgroups = UOp.special((M // block_size) * (N // block_size), "gidx0")
+  sink_inputs = (C.base, A.base, B.base) + tuple(s.base for s in scales) + (threads, workgroups)
+  sink = UOp.sink(*sink_inputs,
+                  arg=KernelInfo(f"hk_fp8_gemm_{M}_{N}_{K}", estimates=Estimates(ops=2*M*N*K, mem=(M*K+N*K)*A.dtype.itemsize+M*N*C.dtype.itemsize)))
+  kittens_path = pathlib.Path(__file__).parent.parent/"thunder"/"amd"
+  src = (kittens_path/"gemm_fp8.cpp").read_text()
+  lib = HIPCCCompiler("gfx950", [f"-I{(kittens_path/'include').as_posix()}", "-std=c++20", "-DKITTENS_CDNA4", "-ffast-math",
+                                 "-DHIP_ENABLE_WARP_SYNC_BUILTINS", f"-DGEMM_M={M}", f"-DGEMM_N={N}", f"-DGEMM_K={K}",
+                                 f"-DSCALE_MODE={scale_mode}"]).compile_cached(src)
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)), UOp(Ops.SOURCE, arg=src),
+                               UOp(Ops.BINARY, arg=lib)))
+
 counters = {"used":0, "todos":[]}
 def todo(msg:str) -> bool: counters["todos"].append(msg); return False
 def _asm_gemm_report():
@@ -2634,7 +2660,7 @@ atexit.register(_asm_gemm_report)

 def can_use_asm_gemm(a:Tensor, b:Tensor) -> bool:
  if a.dtype != b.dtype: return todo(f"dtypes must match {a.dtype} != {b.dtype}")
-  if a.dtype not in {dtypes.bfloat16, dtypes.float16}: return todo(f"only bfloat16/float16, got {a.dtype}")
+  if a.dtype not in {dtypes.bfloat16, dtypes.float16, FP8_DTYPE}: return todo(f"only bfloat16/float16/fp8, got {a.dtype}")
  batch, M, K = (1, *a.shape) if a.ndim == 2 else a.shape
  N = b.shape[1]
  if isinstance(a.device, tuple):
@@ -2647,7 +2673,7 @@ def can_use_asm_gemm(a:Tensor, b:Tensor) -> bool:
    else: return todo(f"sharding mismatch a.ndim={a.ndim} a.uop.axis={a.uop.axis} b.uop.axis={b.uop.axis}")
    dname = a.device[0]
  else: dname = a.device
-  arch = getattr(Device[dname].renderer, "arch", "")
+  arch = Device[dname].renderer.target.arch
  if batch not in {1, 2}: return todo(f"GEMM batch size {batch}")
  # blacklist slow matmul
  # TODO: why is this slow?
@@ -2675,18 +2701,53 @@ def custom_uop_gemm(C:UOp, A:UOp, B:UOp) -> UOp:
 # ** backward gemm, might use the asm gemm

 def custom_gemm_bw(gradient:UOp, kernel:UOp):
-  out, a, b = kernel.src[1:]
-  assert all_same([gradient.device, a.device, b.device, out.device])
-  a_t, b_t, g_t = Tensor(a, device=a.device), Tensor(b, device=a.device), Tensor(gradient, device=a.device)
-  # TODO: this needs to be cleaned up and done properly, the batch dim of grad and a multi need to align
-  g_t = g_t[:a.shape[0]]
-  grad_a = (g_t @ b_t.T).uop
-  grad_b = (a_t.permute(2, 0, 1).reshape(a_t.shape[2], -1) @ g_t.reshape(-1, g_t.shape[-1])).uop
-  return (None, grad_a, grad_b)
+  inputs = kernel.src[1:]
+  if inputs[1].dtype == FP8_DTYPE:
+    grad_amax_state = inputs[5] if len(inputs) == 6 else None
+    out, a, b, s_x, s_w = inputs[:5]
+    a_t, b_t, g_t = Tensor(a, device=a.device), Tensor(b, device=a.device), Tensor(gradient, device=a.device)
+    s_x_t, s_w_t = Tensor(s_x, device=a.device), Tensor(s_w, device=a.device)
+    g_t = g_t[:a.shape[0]]
+    from extra.llama_kernels.cast_amax import _grad_fp8_mailbox
+    from extra.llama_kernels.quantize_fp8_delayed import quantize_fp8_delayed
+    gbase = gradient.base if hasattr(gradient, "base") else gradient
+    mailbox_entry = _grad_fp8_mailbox.pop(gbase, None) or _grad_fp8_mailbox.pop(gradient, None)
+    if mailbox_entry is not None:
+      g_fp8_u, inv_scale_u, _new_amax_u, store_effect = mailbox_entry
+      g_fp8 = Tensor(g_fp8_u, device=a.device)[:a.shape[0]]
+      g_scale = Tensor(inv_scale_u, device=a.device)
+    else:
+      assert grad_amax_state is not None, "fp8 matmul bwd needs either a mailbox entry or a grad_amax_state"
+      g_fp8, g_scale, _, store_effect = quantize_fp8_delayed(g_t, Tensor(grad_amax_state, device=a.device))
+    # dgrad: uses g_scale * x_scale * w_scale
+    grad_a = asm_gemm(g_fp8, b_t, x_scale=g_scale * s_x_t, w_scale=s_w_t)
+    # wgrad: no w_scale
+    g_fp8_2d = g_fp8.reshape(-1, g_fp8.shape[-1])
+    if getenv("FAST_FP8_TRANSPOSE", 0) and g_fp8_2d.shape[0] % 64 == 0 and g_fp8_2d.shape[1] % 64 == 0:
+      from extra.llama_kernels.fp8_transpose import fast_fp8_transpose
+      g_fp8_T = fast_fp8_transpose(g_fp8_2d)
+    else:
+      g_fp8_T = g_fp8.permute(2, 0, 1).reshape(g_t.shape[-1], -1)
+    grad_b = asm_gemm(g_fp8_T, a_t.reshape(-1, a_t.shape[-1]), x_scale=g_scale * s_x_t)
+    # Attach the delayed-amax store effect (if any) to grad_a so realizing grads commits the amax update.
+    ret = (None, grad_a.uop.after(store_effect), grad_b.uop, None, None)
+    if len(inputs) == 6: ret = ret + (None,)
+    return ret
+  else:
+    out, a, b = inputs
+    assert all_same([gradient.device, a.device, b.device, out.device])
+    a_t, b_t, g_t = Tensor(a, device=a.device), Tensor(b, device=a.device), Tensor(gradient, device=a.device)
+    g_t = g_t[:a.shape[0]]
+    if can_use_asm_gemm(g_t, b_t.T): grad_a = asm_gemm(g_t, b_t.T).uop
+    else: grad_a = (g_t @ b_t.T).uop
+    a_t_flat, g_t_flat = a_t.permute(2, 0, 1).reshape(a_t.shape[2], -1), g_t.reshape(-1, g_t.shape[-1])
+    if can_use_asm_gemm(a_t_flat, g_t_flat): grad_b = asm_gemm(a_t_flat, g_t_flat).uop
+    else: grad_b = (a_t_flat @ g_t_flat).uop
+    return (None, grad_a, grad_b)

 # ** main gemm function

-def asm_gemm(a:Tensor, b:Tensor) -> Tensor:
+def asm_gemm(a:Tensor, b:Tensor, x_scale:Tensor|None=None, w_scale:Tensor|None=None, grad_amax_state:Tensor|None=None) -> Tensor:
  assert can_use_asm_gemm(a, b), f"{counters['todos'][-1]}"
  counters["used"] += 1
  unfold_batch = a.ndim == 3 and isinstance(a.device, tuple) and a.uop.axis == 2 and b.uop.axis == 0
@@ -2695,6 +2756,7 @@ def asm_gemm(a:Tensor, b:Tensor) -> Tensor:
    a = a.reshape(a.shape[0]*a.shape[1], a.shape[2])
  squeeze = a.ndim == 2
  if squeeze: a = a.unsqueeze(0)
+  out_dtype = dtypes.bfloat16 if a.dtype == FP8_DTYPE else a.dtype

  batch, M, K = a.shape
  N = b.shape[1]
@@ -2705,19 +2767,27 @@ def asm_gemm(a:Tensor, b:Tensor) -> Tensor:

  if is_multi:
    if n_sharded:
-      out = Tensor(Tensor.invalid(batch, M, N//len(a.device), dtype=a.dtype, device=a.device).uop.multi(2), device=a.device)
+      out = Tensor(Tensor.invalids(batch, M, N//len(a.device), dtype=out_dtype, device=a.device).uop.multi(2), device=a.device)
    elif m_sharded:
-      out = Tensor(Tensor.invalid(batch, M, N, dtype=a.dtype, device=a.device).uop.multi(1), device=a.device)
+      out = Tensor(Tensor.invalids(batch, M, N, dtype=out_dtype, device=a.device).uop.multi(1), device=a.device)
    else:
-      out = Tensor(Tensor.invalid(batch//len(a.device) if a.uop.axis==0 else batch, M, N, dtype=a.dtype, device=a.device).uop.multi(0),
+      out = Tensor(Tensor.invalids(batch//len(a.device) if a.uop.axis==0 else batch, M, N, dtype=out_dtype, device=a.device).uop.multi(0),
                   device=a.device)
  else:
-    out = Tensor.invalid(batch, M, N, dtype=a.dtype, device=a.device)
+    out = Tensor.invalids(batch, M, N, dtype=out_dtype, device=a.device)

-  renderer = Device[a.device[0] if is_multi else a.device].renderer
-  dname, arch = renderer.device, getattr(renderer, "arch", "")
+  renderer = Device[dname:=(a.device[0] if is_multi else a.device)].renderer
+  dname, arch = dname.split(":")[0], renderer.target.arch
  if arch.startswith("gfx950") and getenv("USE_ASM", 1):
-    out = Tensor.custom_kernel(out, a, b, fxn=functools.partial(custom_asm_gemm, dname=dname), grad_fxn=custom_gemm_bw)[0]
+    # fp8 gemm computes [email protected], kernel multiplies output by x_scale * w_scale before bf16 store
+    if a.dtype == FP8_DTYPE:
+      scales = tuple(s for s in (x_scale, w_scale) if s is not None)
+      scale_mode = (1 if x_scale is not None else 0) | (2 if w_scale is not None else 0)
+      extra = [grad_amax_state] if grad_amax_state is not None else []
+      fxn = functools.partial(custom_hk_fp8_gemm, dname=dname, scale_mode=scale_mode)
+      out = Tensor.custom_kernel(out, a, b.T, *scales, *extra, fxn=fxn, grad_fxn=custom_gemm_bw)[0]
+    else:
+      out = Tensor.custom_kernel(out, a, b, fxn=functools.partial(custom_asm_gemm, dname=dname), grad_fxn=custom_gemm_bw)[0]
  else:
    out = Tensor.custom_kernel(out, a, b, fxn=custom_uop_gemm, grad_fxn=custom_gemm_bw)[0]
  if k_sharded: out = out.sum(0)
@@ -1,7 +1,6 @@
 import numpy as np, os
 from tinygrad.helpers import getenv, flat_mv
 from tinygrad import dtypes
-from tinygrad.engine.realize import get_program

 # for copied uops
 from tinygrad import dtypes
@@ -0,0 +1,248 @@
+# RDNA4 128x128 GEMM using WMMA — optimized DS scheduling
+import numpy as np
+from tinygrad import Tensor, Device, Context, GlobalCounters
+from tinygrad.uop.ops import UOp, Ops, KernelInfo
+from tinygrad.helpers import getenv, colored
+from tinygrad.dtype import dtypes, AddrSpace
+from tinygrad.engine.realize import Estimates, run_linear
+from tinygrad.renderer.amd.dsl import s, v, VCC_LO, NULL, src, ttmp
+from tinygrad.runtime.autogen.amd.rdna4.ins import *
+
+BLOCK_M, BLOCK_N, BLOCK_K = 128, 128, 16
+TILES_M, TILES_N = 4, 4
+THREADS, ELEM = 128, 2
+LDS_A_ROW = BLOCK_K*ELEM   # 32
+LDS_B_ROW = BLOCK_N*ELEM  # 256
+LDS_A_SIZE = BLOCK_M * LDS_A_ROW  # 4096
+LDS_B_SIZE = BLOCK_K * LDS_B_ROW  # 4096
+LDS_SIZE = LDS_A_SIZE + LDS_B_SIZE  # 8192
+LDS_B_OFF = LDS_A_SIZE
+ACC, DA, DB, FA, FB, ET = 60, 188, 196, 204, 44, 10
+
+def build_kernel(N, arch='gfx1200'):
+  assert N % BLOCK_M == 0 and N >= 256
+  NO_ALU, NO_DS, NO_GLOBAL = getenv("NO_ALU", 0), getenv("NO_DS", 0), getenv("NO_GLOBAL", 0)
+  I, L, B = [], {}, []
+  def e(i): I.append(i); return i
+  def label(n): L[n] = sum(i.size() for i in I)
+  def br(i, t): B.append((len(I)-1, t))
+
+  e(s_load_b128(sdata=s[4:7], sbase=s[0:1], ioffset=0, soffset=NULL))
+  e(s_load_b64(sdata=s[8:9], sbase=s[0:1], ioffset=0x10, soffset=NULL))
+  e(s_wait_kmcnt(simm16=0))
+  e(s_mov_b32(s[10], ttmp[9])); e(s_and_b32(s[11], ttmp[7], 0xFFFF))
+  e(s_lshl_b32(s[10], s[10], 7)); e(s_lshl_b32(s[11], s[11], 7))
+  e(s_mov_b32(s[12], N)); e(s_lshl_b32(s[13], s[12], 1))
+  e(s_mul_i32(s[14], s[12], BLOCK_K*ELEM))
+  e(s_add_co_i32(s[17], s[12], -2*BLOCK_K))  # loop bound
+
+  e(v_and_b32_e32(v[1], 31, v[0])); e(v_lshrrev_b32_e32(v[2], 5, v[0]))
+  e(v_and_b32_e32(v[3], 1, v[2])); e(v_lshrrev_b32_e32(v[2], 1, v[2]))
+
+  e(v_lshlrev_b32_e32(v[4], 5, v[0]))
+  # B store: transposed layout for stride-32 reads. addr = LDS_B_OFF + (tid%8)*512 + (tid/8)*32
+  e(v_and_b32_e32(v[48], 7, v[0])); e(v_lshlrev_b32_e32(v[5], 9, v[48]))   # (tid%8)*512
+  e(v_lshrrev_b32_e32(v[48], 3, v[0])); e(v_lshlrev_b32_e32(v[48], 5, v[48]))  # (tid/8)*32
+  e(v_add_nc_u32_e32(v[5], v[5], v[48])); e(v_add_nc_u32_e32(v[5], LDS_B_OFF, v[5]))
+
+  e(v_add_nc_u32_e32(v[48], s[11], v[0]))
+  e(v_mul_lo_u32(v[6], v[48], N*ELEM)); e(v_mov_b32_e32(v[7], 0))
+  e(v_lshrrev_b32_e32(v[48], 3, v[0])); e(v_mul_lo_u32(v[8], v[48], N*ELEM))
+  e(v_and_b32_e32(v[48], 7, v[0])); e(v_lshlrev_b32_e32(v[48], 5, v[48]))
+  e(v_add_nc_u32_e32(v[8], v[8], v[48]))
+  e(s_mul_i32(s[15], s[10], ELEM)); e(v_add_nc_u32_e32(v[8], s[15], v[8]))
+  e(v_mov_b32_e32(v[9], 0))
+
+  # LDS read addrs with padded strides (eliminates bank conflicts)
+  # A: (lane%16)*LDS_A_ROW + (lane/16)*16 + wave_m*64*LDS_A_ROW
+  # B: (lane%16)*LDS_B_ROW + (lane/16)*16 + wave_n*64*ELEM + LDS_B_OFF
+  LLA, LLB = 40, 43
+  e(v_and_b32_e32(v[50], 15, v[1])); e(v_lshrrev_b32_e32(v[51], 4, v[1]))
+  e(v_lshlrev_b32_e32(v[LLA], 5, v[50]))        # (lane%16) * 32
+  e(v_lshlrev_b32_e32(v[51], 4, v[51]))        # (lane/16) * 16
+  e(v_add_nc_u32_e32(v[LLA], v[LLA], v[51]))
+  e(v_lshlrev_b32_e32(v[52], 11, v[2]))         # wave_m * 2048
+  e(v_add_nc_u32_e32(v[LLA], v[LLA], v[52]))
+  # B read: transposed layout. addr = LDS_B_OFF + (lane%16)*32 + (lane/16)*16 + wave_n*2*512
+  # wave_n selects column panels: wave_n*2 panels (each panel=16 cols, wave_n covers 64 cols = 4 panels)
+  # But wave_n*2*512 = wave_n*1024. Hmm, wave_n covers cols [wave_n*64 : (wave_n+1)*64].
+  # Each panel = 16 cols = 512 bytes. wave_n*64/16 = wave_n*4 panels. Offset = wave_n*4*512 = wave_n*2048.
+  e(v_lshlrev_b32_e32(v[LLB], 5, v[50]))         # (lane%16) * 32 (stride 32!)
+  e(v_add_nc_u32_e32(v[LLB], v[LLB], v[51]))     # + (lane/16)*16
+  e(v_lshlrev_b32_e32(v[52], 11, v[3]))           # wave_n * 2048
+  e(v_add_nc_u32_e32(v[LLB], v[LLB], v[52]))
+  e(v_add_nc_u32_e32(v[LLB], LDS_B_OFF, v[LLB]))
+
+  for i in range(0, 128, 2):
+    e(VOPD(VOPDOp.V_DUAL_MOV_B32, VOPDOp.V_DUAL_MOV_B32, vdstx=v[ACC+i], vdsty=v[ACC+i+1], srcx0=0, srcy0=0))
+  e(s_mov_b32(s[16], 0))
+
+  if not NO_GLOBAL:
+    for i in range(2): e(global_load_b128(vdst=v[DA+i*4:DA+i*4+3], vaddr=v[6:7], saddr=s[4:5], ioffset=i*16))
+    for i in range(2): e(global_load_b128(vdst=v[DB+i*4:DB+i*4+3], vaddr=v[8:9], saddr=s[6:7], ioffset=i*16))
+    e(s_wait_loadcnt(simm16=0))
+  if not NO_DS:
+    for i in range(2): e(ds_store_b128(addr=v[4], data0=v[DA+i*4:DA+i*4+3], offset0=(i*16)&0xFF, offset1=(i*16)>>8))
+    for i in range(2): e(ds_store_b128(addr=v[5], data0=v[DB+i*4:DB+i*4+3], offset0=(i*16)&0xFF, offset1=(i*16)>>8))
+  if not NO_GLOBAL:
+    e(v_add_nc_u32_e32(v[6], BLOCK_K*ELEM, v[6]))
+    e(v_add_nc_u32_e32(v[8], s[14], v[8]))
+
+  # =============================================================================
+  def emit_iter_body(load_set='AB'):
+    if not NO_DS:
+      e(s_wait_dscnt(simm16=0))
+      e(s_barrier_signal(ssrc0=src[193])); e(s_barrier_wait(simm16=0xFFFF))
+    if not NO_GLOBAL:
+      if 'A' in load_set:
+        for i in range(2): e(global_load_b128(vdst=v[DA+i*4:DA+i*4+3], vaddr=v[6:7], saddr=s[4:5], ioffset=i*16))
+        e(v_add_nc_u32_e32(v[6], BLOCK_K*ELEM, v[6]))
+      if 'B' in load_set:
+        for i in range(2): e(global_load_b128(vdst=v[DB+i*4:DB+i*4+3], vaddr=v[8:9], saddr=s[6:7], ioffset=i*16))
+        e(v_add_nc_u32_e32(v[8], s[14], v[8]))
+    if not NO_DS:
+      # Issue 6 loads: A[0:3] + B[0] + B[1]. B[2:3] interleaved with WMMAs.
+      for tm in range(TILES_M):
+        aoff = tm * 16 * LDS_A_ROW
+        e(ds_load_b128(vdst=v[FA+tm*4:FA+tm*4+3], addr=v[LLA], offset0=aoff&0xFF, offset1=aoff>>8))
+      e(ds_load_b128(vdst=v[FB:FB+3], addr=v[LLB], offset0=0, offset1=0))
+      e(ds_load_b128(vdst=v[FB+4:FB+7], addr=v[LLB], offset0=0, offset1=2))
+      e(s_wait_dscnt(simm16=0))  # wait for 6 loads (no stall!)
+    if not NO_ALU:
+      # B[0] WMMAs — issue B[2] during compute
+      if not NO_DS: e(ds_load_b128(vdst=v[FB+8:FB+11], addr=v[LLB], offset0=0, offset1=4))
+      for tm in range(TILES_M):
+        ac = ACC + (tm*TILES_N+0)*8
+        e(v_wmma_f32_16x16x16_f16(vdst=v[ac:ac+7], src0=v[FA+tm*4:FA+tm*4+3], src1=v[FB:FB+3], src2=v[ac:ac+7]))
+      # B[1] WMMAs — issue B[3] during compute
+      if not NO_DS:
+        e(ds_load_b128(vdst=v[FB+12:FB+15], addr=v[LLB], offset0=0, offset1=6))
+      for tm in range(TILES_M):
+        ac = ACC + (tm*TILES_N+1)*8
+        e(v_wmma_f32_16x16x16_f16(vdst=v[ac:ac+7], src0=v[FA+tm*4:FA+tm*4+3], src1=v[FB+4:FB+7], src2=v[ac:ac+7]))
+      # B[2] WMMAs — B[2] loaded during B[0] WMMAs (~100 cycles ago)
+      if not NO_DS: e(s_wait_dscnt(simm16=1))  # B[2] done, B[3] may still be loading
+      for tm in range(TILES_M):
+        ac = ACC + (tm*TILES_N+2)*8
+        e(v_wmma_f32_16x16x16_f16(vdst=v[ac:ac+7], src0=v[FA+tm*4:FA+tm*4+3], src1=v[FB+8:FB+11], src2=v[ac:ac+7]))
+      # B[3] WMMAs
+      if not NO_DS: e(s_wait_dscnt(simm16=0))
+      for tm in range(TILES_M):
+        ac = ACC + (tm*TILES_N+3)*8
+        e(v_wmma_f32_16x16x16_f16(vdst=v[ac:ac+7], src0=v[FA+tm*4:FA+tm*4+3], src1=v[FB+12:FB+15], src2=v[ac:ac+7]))
+    if not NO_GLOBAL and not NO_DS: e(s_wait_loadcnt(simm16=0))
+    if not NO_DS:
+      for i in range(2): e(ds_store_b128(addr=v[4], data0=v[DA+i*4:DA+i*4+3], offset0=(i*16)&0xFF, offset1=(i*16)>>8))
+      for i in range(2): e(ds_store_b128(addr=v[5], data0=v[DB+i*4:DB+i*4+3], offset0=(i*16)&0xFF, offset1=(i*16)>>8))
+    e(s_add_co_i32(s[16], s[16], BLOCK_K))
+
+  label('LOOP')
+  emit_iter_body(load_set='A')
+  emit_iter_body(load_set='B')
+  e(s_cmp_lt_i32(s[16], s[17])); e(s_cbranch_scc1(simm16=0)); br(I[-1], 'LOOP')
+
+  emit_iter_body(load_set='AB')  # tail with prefetch
+
+  # Final iteration: no prefetch, no ds_store needed
+  if not NO_DS:
+    e(s_wait_dscnt(simm16=0))
+    e(s_barrier_signal(ssrc0=src[193])); e(s_barrier_wait(simm16=0xFFFF))
+  if not NO_DS:
+    for tm in range(TILES_M):
+      aoff = tm * 16 * LDS_A_ROW
+      e(ds_load_b128(vdst=v[FA+tm*4:FA+tm*4+3], addr=v[LLA], offset0=aoff&0xFF, offset1=aoff>>8))
+    e(ds_load_b128(vdst=v[FB:FB+3], addr=v[LLB], offset0=0, offset1=0))
+    e(ds_load_b128(vdst=v[FB+4:FB+7], addr=v[LLB], offset0=0, offset1=2))
+    e(s_wait_dscnt(simm16=0))
+  if not NO_ALU:
+    if not NO_DS: e(ds_load_b128(vdst=v[FB+8:FB+11], addr=v[LLB], offset0=0, offset1=4))
+    for tm in range(TILES_M):
+      ac = ACC + (tm*TILES_N+0)*8
+      e(v_wmma_f32_16x16x16_f16(vdst=v[ac:ac+7], src0=v[FA+tm*4:FA+tm*4+3], src1=v[FB:FB+3], src2=v[ac:ac+7]))
+    if not NO_DS: e(ds_load_b128(vdst=v[FB+12:FB+15], addr=v[LLB], offset0=0, offset1=6))
+    for tm in range(TILES_M):
+      ac = ACC + (tm*TILES_N+1)*8
+      e(v_wmma_f32_16x16x16_f16(vdst=v[ac:ac+7], src0=v[FA+tm*4:FA+tm*4+3], src1=v[FB+4:FB+7], src2=v[ac:ac+7]))
+    if not NO_DS: e(s_wait_dscnt(simm16=1))
+    for tm in range(TILES_M):
+      ac = ACC + (tm*TILES_N+2)*8
+      e(v_wmma_f32_16x16x16_f16(vdst=v[ac:ac+7], src0=v[FA+tm*4:FA+tm*4+3], src1=v[FB+8:FB+11], src2=v[ac:ac+7]))
+    if not NO_DS: e(s_wait_dscnt(simm16=0))
+    for tm in range(TILES_M):
+      ac = ACC + (tm*TILES_N+3)*8
+      e(v_wmma_f32_16x16x16_f16(vdst=v[ac:ac+7], src0=v[FA+tm*4:FA+tm*4+3], src1=v[FB+12:FB+15], src2=v[ac:ac+7]))
+
+  label('EPILOGUE')
+  e(v_and_b32_e32(v[ET], 15, v[1]))
+  e(v_lshrrev_b32_e32(v[ET+1], 4, v[1])); e(v_lshlrev_b32_e32(v[ET+1], 3, v[ET+1]))
+  e(v_lshlrev_b32_e32(v[ET+2], 6, v[2])); e(v_add_nc_u32_e32(v[ET+2], s[11], v[ET+2]))
+  e(v_lshlrev_b32_e32(v[ET+3], 6, v[3])); e(v_add_nc_u32_e32(v[ET+3], s[10], v[ET+3]))
+  e(v_add_nc_u32_e32(v[ET+3], v[ET+3], v[ET])); e(v_mov_b32_e32(v[ET+5], 0))
+
+  for tm in range(TILES_M):
+    for tn in range(TILES_N):
+      ac = ACC + (tm*TILES_N+tn)*8; r_off, c_off = tm*16, tn*16
+      e(v_add_nc_u32_e32(v[ET+6], r_off, v[ET+2])); e(v_add_nc_u32_e32(v[ET+6], v[ET+1], v[ET+6]))
+      e(v_mul_lo_u32(v[ET+4], v[ET+6], s[12])); e(v_add_nc_u32_e32(v[ET+4], v[ET+4], v[ET+3]))
+      if c_off: e(v_add_nc_u32_e32(v[ET+4], c_off, v[ET+4]))
+      e(v_lshlrev_b32_e32(v[ET+4], 1, v[ET+4]))
+      for elem in range(8):
+        e(v_cvt_f16_f32_e32(v[ET+7], v[ac+elem]))
+        e(global_store_b16(vaddr=v[ET+4:ET+5], vsrc=v[ET+7], saddr=s[8:9]))
+        if elem < 7: e(v_add_nc_u32_e32(v[ET+4], s[13], v[ET+4]))
+
+  e(s_wait_storecnt(simm16=0)); e(s_sendmsg(simm16=3)); e(s_endpgm())
+
+  for idx, target in B:
+    off = (L[target] - sum(i.size() for i in I[:idx+1])) // 4
+    assert -32768 <= off <= 32767; I[idx].simm16 = off
+  return I
+
+N = getenv("N", 4096)
+
+def test_matmul():
+  dev = Device[Device.DEFAULT]
+  arch = getattr(dev.renderer, 'arch', 'gfx1200')
+  print(f"Device arch: {arch}")
+  insts = build_kernel(N, arch)
+
+  rng = np.random.default_rng(42)
+  a = Tensor(rng.random((N, N), dtype=np.float32).astype(np.float16))
+  b = Tensor(rng.random((N, N), dtype=np.float32).astype(np.float16))
+  c = Tensor.empty(N, N, dtype=dtypes.half)
+  Tensor.realize(a, b, c)
+
+  grid, local = (N//BLOCK_N, N//BLOCK_M, 1), (THREADS, 1, 1)
+  print(f"Grid: {grid}, Local: {local}")
+
+  dname = Device.DEFAULT
+  def asm_kernel(A, B, C):
+    gidxs = [UOp.special(n, f"gidx{i}") for i,n in enumerate(grid)]
+    lidxs = [UOp.special(THREADS, "lidx0")]
+    lds = UOp(Ops.DEFINE_LOCAL, dtypes.uint8.ptr(size=max(LDS_SIZE, 65536//getenv("LIMIT_OCC",2)), addrspace=AddrSpace.LOCAL), (), 'lds')
+    sink = UOp.sink(A.base, B.base, C.base, lds, *gidxs, *lidxs,
+                    arg=KernelInfo(name=colored("kernel","cyan"), estimates=Estimates(ops=N*N*N*2, mem=N*N*2*3)))
+    return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=tuple([UOp(Ops.INS, arg=x) for x in insts]))))
+
+  c = Tensor.custom_kernel(a, b, c, fxn=asm_kernel)[2]
+  linear = c.schedule_linear()
+
+  ets = []
+  with Context(DEBUG=2):
+    for _ in range(getenv("CNT", 5)):
+      start = GlobalCounters.time_sum_s
+      run_linear(linear)
+      ets.append(GlobalCounters.time_sum_s - start)
+  print(f"REAL TFLOPS {N*N*N*2 / min(ets) * 1e-12:.2f}")
+
+  if getenv("VERIFY", 1):
+    GlobalCounters.reset()
+    c_np = c.float().numpy()
+    a_np, b_np = a.float().numpy(), b.float().numpy()
+    ref = a_np @ b_np
+    err = np.sqrt(np.mean((c_np - ref)**2)) / np.sqrt(np.mean(ref**2))
+    print(f"relative RMSE {err:.6f}")
+    if err != err or err > 0.05: raise RuntimeError(f"matmul is wrong! RMSE={err}")
+
+if __name__ == "__main__":
+  test_matmul()
@@ -2,6 +2,7 @@ import numpy as np
 from tinygrad import dtypes, Tensor
 from tinygrad.helpers import getenv, get_single_element
 from tinygrad.dtype import _to_np_dtype
+from tinygrad.engine.realize import compile_linear
 from tinygrad.codegen.opt import OptOps

 dtype_in = (dtypes.half if getenv("HALF") else dtypes.bfloat16 if getenv("BFLOAT16") else
@@ -38,10 +39,10 @@ if __name__ == "__main__":
    c = a.matmul(b, dtype=acc_dtype).realize()

  if getenv("SHOULD_USE_TC"):
-    sched = a.matmul(b, dtype=acc_dtype).schedule()
-    ei = get_single_element(sched)
-    ei.lower()
-    assert any(opt.op is OptOps.TC for opt in ei.prg.p.applied_opts), f"TC not triggered, {ei.prg.p.applied_opts}"
+    linear = compile_linear(a.matmul(b, dtype=acc_dtype).schedule_linear())
+    call = get_single_element(list(linear.src))
+    applied_opts = call.src[0].src[0].arg.applied_opts
+    assert any(opt.op is OptOps.TC for opt in applied_opts), f"TC not triggered, {applied_opts}"

  ref = a.numpy().astype(np.float32) @ b.numpy().astype(np.float32)
  res = c.numpy()
@@ -1,7 +1,7 @@
-from tinygrad import Tensor, dtypes, Device
-from tinygrad.helpers import getenv, DEBUG
-from tinygrad.codegen.opt.kernel import Kernel, Opt, OptOps
-from tinygrad.engine.realize import CompiledRunner, ExecItem, get_program
+from tinygrad import Tensor, dtypes, Context
+from tinygrad.helpers import getenv
+from tinygrad.codegen.opt import Opt, OptOps
+from tinygrad.engine.realize import run_linear
 from dataclasses import replace

 N = 4096
@@ -11,9 +11,6 @@ if __name__ == "__main__":
  else:
    A, B = Tensor.empty(N, N, dtype=dtypes.float16), Tensor.empty(N, N, dtype=dtypes.float16)
  C = A.matmul(B)
-  si = C.schedule()[-1]
-  ast = si.ast
-  k = Kernel(ast, opts=Device[Device.DEFAULT].renderer)
  if getenv("GEMV"):
    opts = [
      Opt(op=OptOps.UNROLL, axis=0, amt=8),
@@ -28,10 +25,10 @@ if __name__ == "__main__":
      Opt(op=OptOps.LOCAL, axis=1, amt=2),
      Opt(op=OptOps.LOCAL, axis=0, amt=2),
    ]
-  k.apply_opts(opts)
-  prg = get_program(k.ast, k.opts, k.applied_opts)
-  new_src = prg.src
-  # can mod source here
-  prg = replace(prg, src=new_src)
-  ei = ExecItem(si.ast, [x.ensure_allocated() for x in si.bufs], si.metadata, prg=CompiledRunner(prg))
-  for i in range(5): ei.run(wait=True)
+  linear = C.schedule_linear()
+  call = linear.src[-1]
+  new_ast = call.src[0].replace(arg=replace(call.src[0].arg, opts_to_apply=tuple(opts)))
+  new_call = call.replace(src=(new_ast, *call.src[1:]))
+  linear = linear.replace(src=tuple(new_call if c is call else c for c in linear.src))
+  with Context(DEBUG=2):
+    for i in range(5): run_linear(linear)
@@ -4,7 +4,9 @@ import triton.language as tl
 from triton.compiler import AttrsDescriptor, ASTSource, compile as triton_compile
 import numpy as np
 from tinygrad import Tensor, dtypes, Device
-from tinygrad.engine.realize import CompiledRunner, ExecItem, ProgramSpec
+from tinygrad.engine.realize import get_runtime
+from tinygrad.codegen import to_program
+from tinygrad.uop.ops import Ops, UOp, KernelInfo, ProgramInfo
 from tinygrad.helpers import getenv
 np.set_printoptions(suppress=True)

@@ -73,8 +75,11 @@ if __name__ == "__main__":

  A, B = Tensor.normal(M, K, std=1e-1, dtype=dtypes.float16).realize(), Tensor.normal(K, N, std=1e-1, dtype=dtypes.float16).realize()
  C = A.matmul(B)
-  sched = C.schedule()
-  si = sched[-1]
+  from tinygrad.uop.ops import Ops
+  linear, var_vals = C.linear_with_vars()
+  last_call = linear.src[-1]
+  ast = last_call.src[0]
+  bufs = [s.buffer for s in last_call.src[1:] if s.op is not Ops.BIND]

  src = compiled.asm["ptx"]
  # specify the shared memory here so we don't need to do it dynamically
@@ -85,22 +90,27 @@ if __name__ == "__main__":
  # remove debug sections
  src = src.split("\t.file")[0]
  assert '.extern .shared' not in src
-  prg = ProgramSpec("matmul_kernel", src, device=Device.DEFAULT,
-                global_size=[M//BLOCK_SIZE_M, N//BLOCK_SIZE_N, 1], local_size=[32*compiled.metadata.num_warps, 1, 1],
-                mem_estimate=A.nbytes() + B.nbytes() + C.nbytes())
-  ei = ExecItem(si.ast, [x.ensure_allocated() for x in si.bufs], si.metadata, prg=CompiledRunner(prg))
+  info = ProgramInfo(name="matmul_kernel",
+                     global_size=(M//BLOCK_SIZE_M, N//BLOCK_SIZE_N, 1), local_size=(32*compiled.metadata.num_warps, 1, 1))
+  sink = UOp.sink(arg=KernelInfo(name="matmul_kernel"))
+  prg_uop = to_program(UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=Device.DEFAULT), UOp(Ops.LINEAR), UOp(Ops.SOURCE, arg=src)), arg=info),
+                       Device.default.renderer)
+  rt = get_runtime(Device.DEFAULT, prg_uop)
+  all_bufs = [x.ensure_allocated() for x in bufs]
+  prg_bufs = [all_bufs[i] for i in info.globals]
+  gsize, lsize = info.launch_dims({})
  tflops = []
  for i in range(5):
-    tm = ei.run(wait=True)
+    tm = rt(*[b._buf for b in prg_bufs], global_size=gsize, local_size=lsize, vals=info.vals({}), wait=True)
    tflops.append((2*M*K*N/tm)*1e-12)
  print(f"TFLOPS: {max(tflops):.2f}")

  # check correctness
  if getenv("VERIFY"):
-    from tinygrad.engine.realize import run_schedule
+    from tinygrad.engine.realize import run_linear
    triton_buf = np.frombuffer(si.bufs[0].as_memoryview(), np.float16).reshape(M,N)
    print(triton_buf)
-    run_schedule(sched)
+    run_linear(linear, var_vals)
    tinygrad_buf = np.frombuffer(si.bufs[0].as_memoryview(), np.float16).reshape(M,N)
    print(tinygrad_buf)
    np.testing.assert_allclose(triton_buf, tinygrad_buf)
@@ -36,10 +36,10 @@ A = Tensor.rand(M, K, device="CPU")
 B = Tensor.rand(K, N, device="CPU")
 C = (A.reshape(M, 1, K) * B.permute(1,0).reshape(1, N, K)).sum(axis=2)

-sched = C.schedule()
+linear = C.schedule_linear()
 from tinygrad.codegen.opt.kernel import Kernel
 from tinygrad.device import CompilerOptions
-lin = Kernel(sched[-1].ast, CompilerOptions(has_local=False, supports_float4=False))
+lin = Kernel(linear.src[-1].src[0], CompilerOptions(has_local=False, supports_float4=False))
 lin.to_program()
 from tinygrad.runtime.ops_cpu import renderer
 src = renderer("mmult", lin.uops)
@@ -0,0 +1,51 @@
+from __future__ import annotations
+import functools, pathlib
+from tinygrad import Tensor, dtypes
+from tinygrad.uop.ops import Ops
+from tinygrad.runtime.support.compiler_amd import HIPCCCompiler
+
+FP8_MAX = 448.0
+NUM_WG, THREADS_PER_WG = 1024, 256
+
+# per-device abs max without allreduce
+@functools.cache
+def _local_abs_max_fxn(x_p, device):
+  x = Tensor(x_p, device=device)
+  inner = Tensor(x.uop.src[0]) if x.uop.op is Ops.MULTI else x
+  return (inner.abs().max(),)
+
+def local_abs_max(x:Tensor) -> Tensor:
+  param = x.as_param(0)
+  fxn = _local_abs_max_fxn(param.uop, x.device)
+  return Tensor(fxn[0].uop.call(x.uop).gettuple(0))
+
+def scalar_amax(amax_buf:Tensor) -> Tensor:
+  if isinstance(amax_buf.device, tuple):
+    return local_abs_max(amax_buf).detach()
+  return amax_buf.max().detach()
+
+def shard_shape(shape:tuple, axis:int, ndev:int) -> list:
+  s = list(shape)
+  s[axis] //= ndev
+  return s
+
+def dname_of(device) -> str:
+  if isinstance(device, tuple): return device[0].split(":")[0]
+  return device.split(":")[0] if isinstance(device, str) else device
+
+def alloc_like(shape, dtype, device, axis=None) -> Tensor:
+  if isinstance(device, tuple) and axis is not None:
+    return Tensor(Tensor.invalids(*shard_shape(shape, axis, len(device)), dtype=dtype, device=device).uop.multi(axis), device=device)
+  return Tensor.invalids(*shape, dtype=dtype, device=device)
+
+def alloc_local(shape, dtype, device, axis=None) -> Tensor:
+  if isinstance(device, tuple) and axis is not None:
+    return Tensor(Tensor.invalids(*shape, dtype=dtype, device=device).uop.multi(0), device=device)
+  return Tensor.invalids(*shape, dtype=dtype, device=device)
+
+def compile_hip(src:str, defines:list[str]):
+  return HIPCCCompiler("gfx950", ["-std=c++20", "-ffast-math", *defines]).compile_cached(src)
+
+def compile_cpp(cpp_dir:pathlib.Path, cpp_name:str, n_elems:int, hidden:int):
+  src = (cpp_dir/cpp_name).read_text()
+  return src, compile_hip(src, [f"-DN_ELEMS={n_elems}", f"-DHIDDEN={hidden}", f"-DNUM_WG={NUM_WG}", f"-DTHREADS_PER_WG={THREADS_PER_WG}"])
@@ -0,0 +1,74 @@
+from __future__ import annotations
+import functools, pathlib
+from tinygrad import Tensor, dtypes
+from tinygrad.uop.ops import UOp, Ops, KernelInfo
+from tinygrad.renderer import Estimates
+from extra.llama_kernels import FP8_MAX, NUM_WG, THREADS_PER_WG, compile_cpp, alloc_like, alloc_local, scalar_amax, dname_of
+
+# module-level mailbox: grad_xw13 UOp -> (grad_xw13_fp8 UOp, inv_scale UOp, new_amax UOp, store_effect)
+# lets cdna_asm_gemm's bwd reuse the fp8 companion produced by the fused silu_mul bwd kernel
+# instead of doing a redundant bf16 -> fp8 quantize.
+_grad_fp8_mailbox:dict = {}
+
+@functools.cache
+def _custom_fused_bwd_w13(grad_xw13:UOp, grad_xw13_fp8:UOp, grad_amax_buf:UOp,
+                          xw13:UOp, grad_x2:UOp, amax_state:UOp, grad_amax_state:UOp, dname:str) -> UOp:
+  hidden = xw13.shape[2] // 2
+  n_elems = xw13.shape[0] * xw13.shape[1] * hidden
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
+  mem = n_elems * 2 * 5 + n_elems * 2 + NUM_WG * 4 + 4
+  sink = UOp.sink(grad_xw13.base, grad_xw13_fp8.base, grad_amax_buf.base,
+                  xw13.base, grad_x2.base, amax_state.base, grad_amax_state.base, threads, workgroups,
+                  arg=KernelInfo(f"fused_silu_mul_bwd_w13_{n_elems}", estimates=Estimates(ops=10*n_elems, mem=mem)))
+  src, lib = compile_cpp(pathlib.Path(__file__).parent, "cast_amax_bwd_w13.cpp", n_elems, hidden)
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=lib)))
+
+@functools.cache
+def _custom_fused_cast_amax_w13(fp8_out:UOp, amax_buf:UOp, xw13:UOp, amax_state:UOp, grad_amax_state:UOp, dname:str) -> UOp:
+  # NOTE: grad_amax_state is plumbed through as an unused fwd input so the bwd kernel can read it via kernel.src
+  hidden = xw13.shape[2] // 2
+  n_elems = xw13.shape[0] * xw13.shape[1] * hidden
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
+  mem = n_elems * 2 * 2 + n_elems + NUM_WG * 4
+  sink = UOp.sink(fp8_out.base, amax_buf.base, xw13.base, amax_state.base, threads, workgroups,
+                  arg=KernelInfo(f"fused_silu_mul_cast_amax_w13_{n_elems}", estimates=Estimates(ops=5*n_elems, mem=mem)))
+  src, lib = compile_cpp(pathlib.Path(__file__).parent, "cast_amax_fwd_w13.cpp", n_elems, hidden)
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=lib)))
+
+def _fused_quantize_bwd_w13(gradient:UOp, kernel:UOp):
+  _, _, xw13, amax_state, grad_amax_state = kernel.src[1:]
+  device = xw13.device
+  axis = xw13.axis if isinstance(device, tuple) else None
+  grad_xw13     = alloc_like(xw13.shape, dtypes.bfloat16, device, axis)
+  grad_xw13_fp8 = alloc_like(xw13.shape, dtypes.fp8e4m3,  device, axis)
+  grad_amax_buf = alloc_local((NUM_WG,), dtypes.float32,  device, axis)
+  grad_amax_state_t = Tensor(grad_amax_state, device=device)
+  fxn = functools.partial(_custom_fused_bwd_w13, dname=dname_of(device))
+  grad_xw13, grad_xw13_fp8, grad_amax_buf, *_ = Tensor.custom_kernel(
+    grad_xw13, grad_xw13_fp8, grad_amax_buf,
+    Tensor(xw13, device=device), Tensor(gradient, device=device).cast(dtypes.bfloat16),
+    Tensor(amax_state, device=device), grad_amax_state_t, fxn=fxn)
+  inv_scale = (grad_amax_state_t.float() + 1e-8) / FP8_MAX
+  new_grad_amax = scalar_amax(grad_amax_buf)
+  store_effect = grad_amax_state_t.uop.store(new_grad_amax.uop)
+  # Stash fp8 companion + amax store for cdna_asm_gemm's bwd to attach to grad_a.
+  _grad_fp8_mailbox[grad_xw13.uop] = (grad_xw13_fp8.uop, inv_scale.uop, new_grad_amax.uop, store_effect)
+  return (None, None, grad_xw13.uop, None, None)
+
+def fused_quantize_fp8_w13(xw13:Tensor, amax_state:Tensor, fp8_dtype, grad_amax_state:Tensor) -> tuple[Tensor, Tensor, Tensor]:
+  # NOTE: silu(xw1)*xw3 -> fp8 + amax over fused xw13 layout. Returns (fp8, inv_scale, new_amax)
+  # grad_amax_state: delayed amax for grad_xw13 fp8 quantization in the backward.
+  assert xw13.dtype == dtypes.bfloat16, f"expected bf16, got {xw13.dtype}"
+  MBS, SEQ, H2 = xw13.shape
+  assert H2 % 2 == 0, f"w13 last-axis must be even, got {H2}"
+  HIDDEN = H2 // 2
+  axis = xw13.uop.axis if isinstance(xw13.device, tuple) else None
+  fp8_out  = alloc_like((MBS, SEQ, HIDDEN), fp8_dtype,      xw13.device, axis)
+  amax_buf = alloc_local((NUM_WG,),         dtypes.float32, xw13.device, axis)
+  fxn = functools.partial(_custom_fused_cast_amax_w13, dname=dname_of(xw13.device))
+  fp8_out, amax_buf, *_ = Tensor.custom_kernel(fp8_out, amax_buf, xw13, amax_state, grad_amax_state,
+                                                fxn=fxn, grad_fxn=_fused_quantize_bwd_w13)
+  inv_scale = (amax_state.float() + 1e-8) / FP8_MAX
+  return fp8_out, inv_scale, scalar_amax(amax_buf)
@@ -0,0 +1,98 @@
+#include <hip/hip_runtime.h>
+#include <hip/hip_bf16.h>
+#include <hip/hip_fp8.h>
+
+#ifndef N_ELEMS
+#define N_ELEMS 234881024
+#endif
+#ifndef HIDDEN
+#define HIDDEN 14336
+#endif
+#ifndef NUM_WG
+#define NUM_WG 1024
+#endif
+#ifndef THREADS_PER_WG
+#define THREADS_PER_WG 256
+#endif
+
+constexpr int VEC = 8;
+constexpr float FP8_MAX = 448.0f;
+
+static_assert(N_ELEMS % VEC == 0, "N_ELEMS must be divisible by VEC");
+static_assert(HIDDEN % VEC == 0, "HIDDEN must be divisible by VEC");
+
+// fused silu*mul backward, three outputs in a single HBM pass:
+//   1) bf16 grad_xw13      — consumed by downstream bf16 autograd chain
+//   2) fp8  grad_xw13_fp8  — delayed-scale quantize using grad_amax_state (mailbox to matmul bwd)
+//   3) fp32 grad_amax_buf  — per-WG partial |grad_xw13|, reduced into next step's grad_amax_state
+// grad_amax_state is read for the fp8 scale. The store of new_grad_amax into grad_amax_state's
+// buffer is built in Python as a separate effect and threaded into grad_a via .after(store).
+extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
+fused_silu_mul_bwd_w13(
+    __hip_bfloat16*       __restrict__ grad_xw13_out,        // bf16, 2*N_ELEMS
+    __hip_fp8_storage_t*  __restrict__ grad_xw13_fp8_out,    // fp8,  2*N_ELEMS
+    float*                __restrict__ grad_amax_buf,        // fp32, NUM_WG per-WG partials
+    const __hip_bfloat16* __restrict__ xw13,                 // bf16, 2*N_ELEMS
+    const __hip_bfloat16* __restrict__ grad_x2,              // bf16, N_ELEMS
+    const float*          __restrict__ amax_state,           // fp32 scalar (fwd x2 amax)
+    const float*          __restrict__ grad_amax_state)      // fp32 scalar (delayed grad amax)
+{
+  __shared__ float sdata[THREADS_PER_WG];
+
+  const int tid = threadIdx.x;
+  const int wg  = blockIdx.x;
+  const int gid = wg * THREADS_PER_WG + tid;
+  const int stride_elems = NUM_WG * THREADS_PER_WG * VEC;
+
+  const float scale = FP8_MAX / (static_cast<float>(*amax_state) + 1e-8f);
+  const float g_scale = FP8_MAX / (static_cast<float>(*grad_amax_state) + 1e-8f);
+  float local_max = 0.0f;
+
+  for (int base = gid * VEC; base < N_ELEMS; base += stride_elems) {
+    const int outer = base / HIDDEN;
+    const int inner = base % HIDDEN;
+    const int xw1_off = outer * 2 * HIDDEN + inner;
+    const int xw3_off = xw1_off + HIDDEN;
+
+    float4 x1_raw = *reinterpret_cast<const float4*>(&xw13[xw1_off]);
+    float4 x3_raw = *reinterpret_cast<const float4*>(&xw13[xw3_off]);
+    float4 g_raw  = *reinterpret_cast<const float4*>(&grad_x2[base]);
+
+    const __hip_bfloat16 *x1 = reinterpret_cast<const __hip_bfloat16*>(&x1_raw);
+    const __hip_bfloat16 *x3 = reinterpret_cast<const __hip_bfloat16*>(&x3_raw);
+    const __hip_bfloat16 *gv = reinterpret_cast<const __hip_bfloat16*>(&g_raw);
+
+    __hip_bfloat16 out1[VEC], out3[VEC];
+    __hip_fp8_storage_t fp8_1[VEC], fp8_3[VEC];
+    #pragma unroll
+    for (int i = 0; i < VEC; i++) {
+      const float f1 = static_cast<float>(x1[i]);
+      const float f3 = static_cast<float>(x3[i]);
+      const float fg = static_cast<float>(gv[i]);
+      const float sig = 1.0f / (1.0f + __expf(-f1));
+      const float silu = f1 * sig;
+      const float silu_prime = sig + silu * (1.0f - sig);
+      const float gs = fg * scale;
+      const float g1 = gs * silu_prime * f3;
+      const float g3 = gs * silu;
+      out1[i] = static_cast<__hip_bfloat16>(g1);
+      out3[i] = static_cast<__hip_bfloat16>(g3);
+      local_max = fmaxf(local_max, fmaxf(fabsf(g1), fabsf(g3)));
+      fp8_1[i] = __hip_cvt_float_to_fp8(fmaxf(-FP8_MAX, fminf(FP8_MAX, g1 * g_scale)), __HIP_SATFINITE, __HIP_E4M3);
+      fp8_3[i] = __hip_cvt_float_to_fp8(fmaxf(-FP8_MAX, fminf(FP8_MAX, g3 * g_scale)), __HIP_SATFINITE, __HIP_E4M3);
+    }
+
+    *reinterpret_cast<float4*>(&grad_xw13_out[xw1_off]) = *reinterpret_cast<float4*>(out1);
+    *reinterpret_cast<float4*>(&grad_xw13_out[xw3_off]) = *reinterpret_cast<float4*>(out3);
+    *reinterpret_cast<uint64_t*>(&grad_xw13_fp8_out[xw1_off]) = *reinterpret_cast<uint64_t*>(fp8_1);
+    *reinterpret_cast<uint64_t*>(&grad_xw13_fp8_out[xw3_off]) = *reinterpret_cast<uint64_t*>(fp8_3);
+  }
+
+  sdata[tid] = local_max;
+  __syncthreads();
+  for (int s = THREADS_PER_WG / 2; s > 0; s >>= 1) {
+    if (tid < s) sdata[tid] = fmaxf(sdata[tid], sdata[tid + s]);
+    __syncthreads();
+  }
+  if (tid == 0) grad_amax_buf[wg] = sdata[0];
+}
@@ -0,0 +1,79 @@
+#include <hip/hip_runtime.h>
+#include <hip/hip_bf16.h>
+#include <hip/hip_fp8.h>
+
+#ifndef N_ELEMS
+#define N_ELEMS 234881024
+#endif
+#ifndef HIDDEN
+#define HIDDEN 14336
+#endif
+#ifndef NUM_WG
+#define NUM_WG 1024
+#endif
+#ifndef THREADS_PER_WG
+#define THREADS_PER_WG 256
+#endif
+
+constexpr int VEC = 8;
+constexpr float FP8_MAX = 448.0f;
+
+static_assert(N_ELEMS % VEC == 0, "N_ELEMS must be divisible by VEC");
+static_assert(HIDDEN % VEC == 0, "HIDDEN must be divisible by VEC (so VEC loads don't straddle block boundary)");
+
+extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
+fused_silu_mul_cast_amax_w13(
+    __hip_fp8_storage_t*  __restrict__ fp8_out,         // fp8, N_ELEMS
+    float*                __restrict__ amax_buf,        // fp32, NUM_WG (per-WG amaxes)
+    const __hip_bfloat16* __restrict__ xw13,            // bf16, 2*N_ELEMS
+    const float*          __restrict__ amax_state)      // fp32 scalar
+{
+  __shared__ float sdata[THREADS_PER_WG];
+
+  const int tid = threadIdx.x;
+  const int wg  = blockIdx.x;
+  const int gid = wg * THREADS_PER_WG + tid;
+  const int stride_elems = NUM_WG * THREADS_PER_WG * VEC;
+
+  const float scale = FP8_MAX / (static_cast<float>(*amax_state) + 1e-8f);
+  float local_max = 0.0f;
+
+  // grid-stride over 8-element groups
+  for (int base = gid * VEC; base < N_ELEMS; base += stride_elems) {
+    // interleaved xw13 layout: xw1 and xw3 are not contiguous halves
+    const int outer = base / HIDDEN;
+    const int inner = base % HIDDEN;
+    const int xw1_off = outer * 2 * HIDDEN + inner;
+    const int xw3_off = xw1_off + HIDDEN;
+
+    float4 x1_raw = *reinterpret_cast<const float4*>(&xw13[xw1_off]);
+    float4 x3_raw = *reinterpret_cast<const float4*>(&xw13[xw3_off]);
+
+    const __hip_bfloat16 *x1 = reinterpret_cast<const __hip_bfloat16*>(&x1_raw);
+    const __hip_bfloat16 *x3 = reinterpret_cast<const __hip_bfloat16*>(&x3_raw);
+
+    __hip_fp8_storage_t out[VEC];
+    #pragma unroll
+    for (int i = 0; i < VEC; i++) {
+      const float f1 = static_cast<float>(x1[i]);
+      const float f3 = static_cast<float>(x3[i]);
+      const float silu = f1 / (1.0f + __expf(-f1));
+      const float x2 = silu * f3;
+      local_max = fmaxf(local_max, fabsf(x2));
+      const float x_scaled = fmaxf(-FP8_MAX, fminf(FP8_MAX, x2 * scale));
+      out[i] = __hip_cvt_float_to_fp8(x_scaled, __HIP_SATFINITE, __HIP_E4M3);
+    }
+
+    *reinterpret_cast<uint64_t*>(&fp8_out[base]) = *reinterpret_cast<uint64_t*>(out);
+  }
+
+  // LDS tree reduction: per-workgroup amax
+  sdata[tid] = local_max;
+  __syncthreads();
+  for (int s = THREADS_PER_WG / 2; s > 0; s >>= 1) {
+    if (tid < s) sdata[tid] = fmaxf(sdata[tid], sdata[tid + s]);
+    __syncthreads();
+  }
+
+  if (tid == 0) amax_buf[wg] = sdata[0];
+}
@@ -0,0 +1,41 @@
+from __future__ import annotations
+import functools, pathlib
+from tinygrad import Tensor, dtypes
+from tinygrad.uop.ops import UOp, Ops, KernelInfo
+from tinygrad.renderer import Estimates
+from extra.llama_kernels import THREADS_PER_WG, alloc_like, dname_of, compile_hip
+
+TILE = 64
+
+@functools.cache
+def _custom_fp8_transpose(out:UOp, inp:UOp, dname:str) -> UOp:
+  M, N = inp.shape
+  num_wg = (M // TILE) * (N // TILE)
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(num_wg, "gidx0")
+  mem = M * N * 2  # one byte read + one byte write per element
+  sink = UOp.sink(out.base, inp.base, threads, workgroups,
+                  arg=KernelInfo(f"fp8_transpose_{M}_{N}",
+                                 estimates=Estimates(ops=M*N, mem=mem)))
+  src = (pathlib.Path(__file__).parent/"fp8_transpose.cpp").read_text()
+  defines = [f"-DM_DIM={M}", f"-DN_DIM={N}", f"-DTHREADS_PER_WG={THREADS_PER_WG}"]
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=compile_hip(src, defines))))
+
+def fast_fp8_transpose(t:Tensor) -> Tensor:
+  assert t.ndim == 2, f"fast_fp8_transpose needs 2D input, got shape {t.shape}"
+  assert t.dtype in dtypes.fp8s, f"fast_fp8_transpose needs fp8 dtype, got {t.dtype}"
+  M, N = t.shape
+  assert M % TILE == 0 and N % TILE == 0, f"M={M}, N={N} must be multiples of {TILE}"
+
+  device = t.device
+  axis = t.uop.axis if isinstance(device, tuple) else None
+  out_axis = None
+  if axis == 0: out_axis = 1
+  elif axis == 1: out_axis = 0
+  elif axis is not None:
+    raise ValueError(f"fast_fp8_transpose: unsupported axis {axis}")
+
+  out = alloc_like((N, M), t.dtype, device, out_axis)
+  fxn = functools.partial(_custom_fp8_transpose, dname=dname_of(device))
+  out, _ = Tensor.custom_kernel(out, t, fxn=fxn)
+  return out
@@ -0,0 +1,74 @@
+#include <hip/hip_runtime.h>
+
+// LDS-staged 64x64 fp8 transpose.
+//   in : (M_DIM, N_DIM) fp8 contiguous
+//   out: (N_DIM, M_DIM) fp8 contiguous, out[c][r] = in[r][c]
+//
+// One WG processes one 64x64 output tile.  Each thread reads one uint4 (16 fp8) coalesced
+// from input rows, stages into LDS, then writes one uint4 coalesced to the output (whose
+// 16 fp8 come from 16 different input rows via in-LDS gather).
+//
+// LDS layout: lds[64][LDS_STRIDE] with LDS_STRIDE=65 (1 byte pad) to mitigate bank conflicts
+// during the column-direction read of the write phase.
+
+#ifndef M_DIM
+#define M_DIM 16384
+#endif
+#ifndef N_DIM
+#define N_DIM 28672
+#endif
+#ifndef THREADS_PER_WG
+#define THREADS_PER_WG 256
+#endif
+
+constexpr int TILE = 64;
+constexpr int VEC = 16;                    // fp8 per uint4 (128-bit) load/store
+constexpr int LDS_PAD = 1;
+constexpr int LDS_STRIDE = TILE + LDS_PAD; // 65 fp8 per row
+
+static_assert(THREADS_PER_WG * VEC == TILE * TILE, "256 threads * 16 fp8 = 64*64");
+static_assert(M_DIM % TILE == 0, "M_DIM must be a multiple of 64");
+static_assert(N_DIM % TILE == 0, "N_DIM must be a multiple of 64");
+
+constexpr int N_TILES_N = N_DIM / TILE;
+
+struct alignas(16) fp8x16 { uint8_t v[16]; };
+
+extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
+fp8_transpose(uint8_t* __restrict__ out,         // (N_DIM, M_DIM)
+              const uint8_t* __restrict__ in)    // (M_DIM, N_DIM)
+{
+  __shared__ uint8_t lds[TILE * LDS_STRIDE];
+
+  const int tid     = threadIdx.x;
+  const int wg_id   = blockIdx.x;
+  const int tile_r  = wg_id / N_TILES_N;       // tile index along M dim of input
+  const int tile_c  = wg_id % N_TILES_N;       // tile index along N dim of input
+
+  const int a = tid / (TILE / VEC);            // 0..63 (row within tile during read; col within tile during write)
+  const int b = tid % (TILE / VEC);            // 0..3
+  const int b16 = b * VEC;                     // 0,16,32,48
+
+  // ---- Read phase: input rows -> LDS rows
+  {
+    const long long src = (long long)(tile_r * TILE + a) * (long long)N_DIM
+                        + (long long)(tile_c * TILE + b16);
+    fp8x16 v = *reinterpret_cast<const fp8x16*>(&in[src]);
+    *reinterpret_cast<fp8x16*>(&lds[a * LDS_STRIDE + b16]) = v;
+  }
+  __syncthreads();
+
+  // ---- Write phase: LDS columns (gathered) -> output rows
+  // out[(tile_c*TILE + a)][(tile_r*TILE + b16 + i)] = in[(tile_r*TILE + b16 + i)][(tile_c*TILE + a)]
+  //                                                  = lds[b16 + i][a]
+  {
+    fp8x16 v;
+    #pragma unroll
+    for (int i = 0; i < VEC; ++i) {
+      v.v[i] = lds[(b16 + i) * LDS_STRIDE + a];
+    }
+    const long long dst = (long long)(tile_c * TILE + a) * (long long)M_DIM
+                        + (long long)(tile_r * TILE + b16);
+    *reinterpret_cast<fp8x16*>(&out[dst]) = v;
+  }
+}
@@ -0,0 +1,96 @@
+from __future__ import annotations
+import functools, pathlib
+from tinygrad import Tensor, dtypes
+from tinygrad.uop.ops import UOp, Ops, KernelInfo
+from tinygrad.renderer import Estimates
+from tinygrad.runtime.support.compiler_amd import HIPCCCompiler
+
+THREADS_PER_WG = 256
+
+@functools.cache
+def _custom_fused_ce_loss_fwd(loss_out:UOp, max_out:UOp, lse_out:UOp, logits:UOp, targets:UOp,
+                              dname:str, vocab:int, rows:int, label_smoothing:float) -> UOp:
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(rows, "gidx0")
+  mem = rows * vocab * 2 + rows * 12 + rows * 4
+  sink = UOp.sink(loss_out.base, max_out.base, lse_out.base, logits.base, targets.base,
+                  threads, workgroups,
+                  arg=KernelInfo(f"fused_ce_loss_fwd", estimates=Estimates(ops=6*rows*vocab, mem=mem)))
+  src = (pathlib.Path(__file__).parent/"fused_ce_loss.cpp").read_text()
+  defines = [f"-DVOCAB={vocab}", f"-DTHREADS_PER_WG={THREADS_PER_WG}",
+             f"-DLABEL_SMOOTHING={label_smoothing}f"]
+  lib = HIPCCCompiler("gfx950", ["-std=c++20", "-ffast-math", *defines]).compile_cached(src)
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=lib)))
+
+@functools.cache
+def _custom_fused_ce_loss_bwd(d_logits:UOp, logits:UOp, lse:UOp, targets:UOp, scale:UOp,
+                              dname:str, vocab:int, rows:int, label_smoothing:float) -> UOp:
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(rows, "gidx0")
+  mem = rows * vocab * 4 + rows * 8 + 4
+  sink = UOp.sink(d_logits.base, logits.base, lse.base, targets.base, scale.base,
+                  threads, workgroups,
+                  arg=KernelInfo(f"fused_ce_loss_bwd", estimates=Estimates(ops=4*rows*vocab, mem=mem)))
+  src = (pathlib.Path(__file__).parent/"fused_ce_loss_bwd.cpp").read_text()
+  defines = [f"-DVOCAB={vocab}", f"-DTHREADS_PER_WG={THREADS_PER_WG}",
+             f"-DLABEL_SMOOTHING={label_smoothing}f"]
+  lib = HIPCCCompiler("gfx950", ["-std=c++20", "-ffast-math", *defines]).compile_cached(src)
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=lib)))
+
+def _fused_ce_loss_bwd(gradient:UOp, kernel:UOp, label_smoothing:float):
+  # NOTE: forward inputs are (loss_out, max_out, lse_out, logits, targets)
+  # gradient is the upstream grad w.r.t. per-row loss (shape: (rows,) fp32)
+  _, _, lse_u, logits_u, targets_u = kernel.src[1:]
+  device = logits_u.device
+  rows, VOCAB = logits_u.shape  # (rows, VOCAB) after reshape
+  if isinstance(device, tuple):
+    axis = logits_u.axis
+    ndev = len(device)
+    d_logits = Tensor(Tensor.invalids(rows // ndev, VOCAB, dtype=dtypes.bfloat16, device=device).uop.multi(axis), device=device)
+    dname = device[0].split(":")[0]
+    rows_per_dev = rows // ndev
+  else:
+    d_logits = Tensor.invalids(rows, VOCAB, dtype=dtypes.bfloat16, device=device)
+    dname = device.split(":")[0] if isinstance(device, str) else device
+    rows_per_dev = rows
+  # NOTE: .mean() backward gives same grad per row (1/N), so broadcast is safe; take scalar
+  scale = Tensor(gradient, device=device).float().reshape(-1)[0:1].contiguous()
+  logits_t = Tensor(logits_u.after(kernel), device=device)
+  lse_t = Tensor(lse_u.after(kernel), device=device)
+  targets_t = Tensor(targets_u, device=device)
+  fxn = functools.partial(_custom_fused_ce_loss_bwd, dname=dname, vocab=VOCAB, rows=rows_per_dev, label_smoothing=label_smoothing)
+  d_logits, *_ = Tensor.custom_kernel(d_logits, logits_t, lse_t, targets_t, scale, fxn=fxn)
+  return (None, None, None, d_logits.uop, None)
+
+def fused_ce_loss(logits:Tensor, targets:Tensor, label_smoothing:float=0.1) -> Tensor:
+  # NOTE: fused sparse_categorical_crossentropy with label smoothing, returns mean loss scalar
+  assert logits.dtype == dtypes.bfloat16, f"expected bf16, got {logits.dtype}"
+  assert logits.ndim == 3, f"expected (MBS, SEQ, VOCAB), got {logits.shape}"
+  MBS, SEQ, VOCAB = logits.shape
+  rows = MBS * SEQ
+  if isinstance(logits.device, tuple):
+    axis = logits.uop.axis
+    assert axis in (0, 1), f"unsupported sharding axis={axis} for CE loss"
+    ndev = len(logits.device)
+    loss_out = Tensor(Tensor.invalids(rows // ndev, dtype=dtypes.float32, device=logits.device).uop.multi(0),
+                      device=logits.device)
+    max_out  = Tensor(Tensor.invalids(rows // ndev, dtype=dtypes.float32, device=logits.device).uop.multi(0),
+                      device=logits.device)
+    lse_out  = Tensor(Tensor.invalids(rows // ndev, dtype=dtypes.float32, device=logits.device).uop.multi(0),
+                      device=logits.device)
+    dname = logits.device[0].split(":")[0]
+    rows_per_dev = rows // ndev
+  else:
+    loss_out = Tensor.invalids(rows, dtype=dtypes.float32, device=logits.device)
+    max_out  = Tensor.invalids(rows, dtype=dtypes.float32, device=logits.device)
+    lse_out  = Tensor.invalids(rows, dtype=dtypes.float32, device=logits.device)
+    dname = logits.device.split(":")[0] if isinstance(logits.device, str) else logits.device
+    rows_per_dev = rows
+  logits_flat = logits.reshape(rows, VOCAB)
+  targets_flat = targets.reshape(-1).cast(dtypes.int32)
+  fxn = functools.partial(_custom_fused_ce_loss_fwd, dname=dname, vocab=VOCAB, rows=rows_per_dev,
+                          label_smoothing=label_smoothing)
+  loss_out, max_out, lse_out, *_ = Tensor.custom_kernel(
+    loss_out, max_out, lse_out, logits_flat, targets_flat,
+    fxn=fxn, grad_fxn=functools.partial(_fused_ce_loss_bwd, label_smoothing=label_smoothing))
+  return loss_out.mean()
@@ -0,0 +1,104 @@
+#include <hip/hip_runtime.h>
+#include <hip/hip_bf16.h>
+
+// Fused forward sparse-CE with label smoothing.
+// SINGLE-PASS online softmax + vectorized 8-wide bf16 loads for HBM coalescing.
+
+#ifndef VOCAB
+#define VOCAB 128256
+#endif
+#ifndef THREADS_PER_WG
+#define THREADS_PER_WG 256
+#endif
+#ifndef LABEL_SMOOTHING
+#define LABEL_SMOOTHING 0.1f
+#endif
+
+constexpr int VEC = 8;
+
+extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
+fused_ce_loss_fwd(
+    float*                __restrict__ loss_out,  // out: fp32, ROWS
+    float*                __restrict__ max_out,   // out: fp32, ROWS
+    float*                __restrict__ lse_out,   // out: fp32, ROWS
+    const __hip_bfloat16* __restrict__ logits,    // in:  bf16, ROWS*VOCAB
+    const int*            __restrict__ targets)   // in:  int32, ROWS
+{
+  __shared__ float sdata_m[THREADS_PER_WG];
+  __shared__ float sdata_s[THREADS_PER_WG];
+  __shared__ float sdata_sumx[THREADS_PER_WG];
+  __shared__ float sdata_tgt[THREADS_PER_WG];
+
+  const int tid = threadIdx.x;
+  const int row = blockIdx.x;
+  const int target = targets[row];
+  const __hip_bfloat16* row_logits = logits + (size_t)row * VOCAB;
+
+  float m = -INFINITY;
+  float s = 0.0f;
+  float sum_x = 0.0f;
+  float target_logit = 0.0f;
+  constexpr bool needs_sum_x = (LABEL_SMOOTHING != 0.0f);
+
+  // Vectorized stride: each iter loads 8 bf16 = 16 bytes. Warp loads 32*16 = 512 bytes (4 cache lines).
+  const int VOCAB_VEC = VOCAB & ~(VEC - 1);  // round down to multiple of VEC
+  for (int i = tid * VEC; i < VOCAB_VEC; i += THREADS_PER_WG * VEC) {
+    float4 raw = *reinterpret_cast<const float4*>(&row_logits[i]);
+    const __hip_bfloat16* xi = reinterpret_cast<const __hip_bfloat16*>(&raw);
+    #pragma unroll
+    for (int k = 0; k < VEC; k++) {
+      const float x = static_cast<float>(xi[k]);
+      if constexpr (needs_sum_x) sum_x += x;
+      if (i + k == target) target_logit = x;
+      if (x > m) {
+        s = s * __expf(m - x) + 1.0f;
+        m = x;
+      } else {
+        s += __expf(x - m);
+      }
+    }
+  }
+  // tail (VOCAB not divisible by VEC):
+  for (int i = VOCAB_VEC + tid; i < VOCAB; i += THREADS_PER_WG) {
+    const float x = static_cast<float>(row_logits[i]);
+    if constexpr (needs_sum_x) sum_x += x;
+    if (i == target) target_logit = x;
+    if (x > m) { s = s * __expf(m - x) + 1.0f; m = x; }
+    else        { s += __expf(x - m); }
+  }
+
+  sdata_m[tid] = m;
+  sdata_s[tid] = s;
+  sdata_sumx[tid] = sum_x;
+  sdata_tgt[tid] = target_logit;
+  __syncthreads();
+
+  for (int step = THREADS_PER_WG / 2; step > 0; step >>= 1) {
+    if (tid < step) {
+      const float m1 = sdata_m[tid];
+      const float m2 = sdata_m[tid + step];
+      const float s1 = sdata_s[tid];
+      const float s2 = sdata_s[tid + step];
+      const float m_new = fmaxf(m1, m2);
+      const float s_new = s1 * __expf(m1 - m_new) + s2 * __expf(m2 - m_new);
+      sdata_m[tid] = m_new;
+      sdata_s[tid] = s_new;
+      sdata_sumx[tid] += sdata_sumx[tid + step];
+      sdata_tgt[tid]  += sdata_tgt[tid + step];
+    }
+    __syncthreads();
+  }
+
+  if (tid == 0) {
+    const float row_max = sdata_m[0];
+    const float row_sum_exp = sdata_s[0];
+    const float row_sum_x = sdata_sumx[0];
+    const float tgt = sdata_tgt[0];
+    const float row_lse = logf(row_sum_exp) + row_max;
+    const float mean_logits = row_sum_x / static_cast<float>(VOCAB);
+    const float loss = row_lse - (1.0f - LABEL_SMOOTHING) * tgt - LABEL_SMOOTHING * mean_logits;
+    loss_out[row] = loss;
+    max_out[row]  = row_max;
+    lse_out[row]  = row_lse;
+  }
+}
@@ -0,0 +1,58 @@
+#include <hip/hip_runtime.h>
+#include <hip/hip_bf16.h>
+
+// Vectorized CE bwd: 8-wide bf16 loads + stores.
+
+#ifndef VOCAB
+#define VOCAB 128256
+#endif
+#ifndef THREADS_PER_WG
+#define THREADS_PER_WG 256
+#endif
+#ifndef LABEL_SMOOTHING
+#define LABEL_SMOOTHING 0.1f
+#endif
+
+constexpr int VEC = 8;
+
+extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
+fused_ce_loss_bwd(
+    __hip_bfloat16*       __restrict__ d_logits,
+    const __hip_bfloat16* __restrict__ logits,
+    const float*          __restrict__ lse,
+    const int*            __restrict__ targets,
+    const float*          __restrict__ scale_in)
+{
+  const int tid = threadIdx.x;
+  const int row = blockIdx.x;
+  const int target = targets[row];
+  const float lse_r = lse[row];
+  const __hip_bfloat16* row_logits  = logits   + (size_t)row * VOCAB;
+  __hip_bfloat16*       row_dlogits = d_logits + (size_t)row * VOCAB;
+  const float inv_vocab = 1.0f / static_cast<float>(VOCAB);
+  const float scale = *scale_in;
+  const float ls_term = LABEL_SMOOTHING * inv_vocab;
+
+  const int VOCAB_VEC = VOCAB & ~(VEC - 1);
+  for (int i = tid * VEC; i < VOCAB_VEC; i += THREADS_PER_WG * VEC) {
+    float4 raw = *reinterpret_cast<const float4*>(&row_logits[i]);
+    const __hip_bfloat16* xi = reinterpret_cast<const __hip_bfloat16*>(&raw);
+    __hip_bfloat16 out[VEC];
+    #pragma unroll
+    for (int k = 0; k < VEC; k++) {
+      const float x = static_cast<float>(xi[k]);
+      float g = __expf(x - lse_r);
+      if (i + k == target) g -= (1.0f - LABEL_SMOOTHING);
+      g -= ls_term;
+      out[k] = static_cast<__hip_bfloat16>(g * scale);
+    }
+    *reinterpret_cast<float4*>(&row_dlogits[i]) = *reinterpret_cast<float4*>(out);
+  }
+  for (int i = VOCAB_VEC + tid; i < VOCAB; i += THREADS_PER_WG) {
+    const float x = static_cast<float>(row_logits[i]);
+    float g = __expf(x - lse_r);
+    if (i == target) g -= (1.0f - LABEL_SMOOTHING);
+    g -= ls_term;
+    row_dlogits[i] = static_cast<__hip_bfloat16>(g * scale);
+  }
+}
@@ -0,0 +1,55 @@
+from __future__ import annotations
+import functools, pathlib
+from tinygrad import Tensor, dtypes
+from tinygrad.uop.ops import UOp, Ops, KernelInfo
+from tinygrad.renderer import Estimates
+from extra.llama_kernels import THREADS_PER_WG, dname_of, compile_hip
+
+ELEMS_PER_THREAD = 8  # vectorized 16-byte load (uint4 = 8 bf16)
+
+def _build_src(n_chunks:int) -> str:
+  template = (pathlib.Path(__file__).parent/"fused_pad_grad_accum.cpp").read_text()
+  params = "".join(f",\n    const __hip_bfloat16* __restrict__ chunk{i}" for i in range(n_chunks))
+  dispatch = "\n    ".join(f"case {i}: chunk_ptr = chunk{i}; break;" for i in range(n_chunks))
+  return (template.replace("__FUSED_PAD_GRAD_ACCUM_PARAMS", params)
+                  .replace("__FUSED_PAD_GRAD_ACCUM_DISPATCH", dispatch))
+
+@functools.cache
+def _custom_fused_pad_grad_accum(grad_buf:UOp, *chunk_uops, dname:str, n_chunks:int, chunk_size:int) -> UOp:
+  total = n_chunks * chunk_size
+  elems_per_block = THREADS_PER_WG * ELEMS_PER_THREAD
+  assert chunk_size % elems_per_block == 0, f"chunk_size {chunk_size} must be multiple of {elems_per_block}"
+  num_wg = total // elems_per_block
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(num_wg, "gidx0")
+  mem = total * 2 * 3
+  sink = UOp.sink(grad_buf.base, *(c.base for c in chunk_uops), threads, workgroups,
+                  arg=KernelInfo(f"fused_pad_grad_accum_n{n_chunks}_c{chunk_size}",
+                                 estimates=Estimates(ops=2*total, mem=mem)))
+  src = _build_src(n_chunks)
+  defines = [f"-DCHUNK_SIZE={chunk_size}", f"-DTHREADS_PER_WG={THREADS_PER_WG}", f"-DELEMS_PER_THREAD={ELEMS_PER_THREAD}"]
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=compile_hip(src, defines))))
+
+def can_fused_pad_grad_accum(grad_buf:Tensor, chunks:list[Tensor]) -> bool:
+  if not chunks or grad_buf.dtype != dtypes.bfloat16: return False
+  if any(c.dtype != dtypes.bfloat16 for c in chunks): return False
+  chunk_shape = chunks[0].shape
+  if any(c.shape != chunk_shape for c in chunks): return False
+  chunk_size, total = 1, 1
+  for d in chunk_shape: chunk_size *= d
+  for d in grad_buf.shape: total *= d
+  return total == len(chunks) * chunk_size and chunk_size % (THREADS_PER_WG * ELEMS_PER_THREAD) == 0
+
+def fused_pad_grad_accum(grad_buf:Tensor, chunks:list[Tensor]) -> Tensor:
+  # NOTE: grad_buf += cat(*chunks, dim=0) in one HBM pass (in-place add). Returns new grad_buf Tensor.
+  # Requires uniform chunk shapes and chunk_size % (THREADS_PER_WG*ELEMS_PER_THREAD) == 0.
+  assert chunks and grad_buf.dtype == dtypes.bfloat16
+  for c in chunks: assert c.dtype == dtypes.bfloat16, f"chunk dtype must be bf16, got {c.dtype}"
+  chunk_size, total = 1, 1
+  for d in chunks[0].shape: chunk_size *= d
+  for d in grad_buf.shape: total *= d
+  assert total == len(chunks) * chunk_size, f"grad_buf size {total} != n_chunks {len(chunks)} * chunk_size {chunk_size}"
+  fxn = functools.partial(_custom_fused_pad_grad_accum, dname=dname_of(grad_buf.device),
+                          n_chunks=len(chunks), chunk_size=chunk_size)
+  out, *_ = Tensor.custom_kernel(grad_buf, *chunks, fxn=fxn)
+  return out
@@ -0,0 +1,63 @@
+// Fused custom kernel: grad_buf += cat(*chunks, dim=0) in one HBM pass.
+//
+// Template source — chunk parameter list and switch dispatch are filled by codegen
+// in cast_amax.py:_build_fused_pad_grad_accum_src to support arbitrary N.
+//
+// Defines required at compile time:
+//   CHUNK_SIZE       elements per chunk (must be multiple of THREADS_PER_WG * ELEMS_PER_THREAD)
+//   THREADS_PER_WG
+//   ELEMS_PER_THREAD (8 = one uint4 per thread = 16-byte vectorized load)
+//
+// Layout: one block-per-(slice-of-chunk) — blockIdx.x / BLOCKS_PER_CHUNK selects the chunk.
+// All threads in a block read the same chunk → switch is uniform → no warp divergence.
+
+#include <hip/hip_runtime.h>
+#include <hip/hip_bf16.h>
+
+#ifndef THREADS_PER_WG
+#define THREADS_PER_WG 256
+#endif
+#ifndef ELEMS_PER_THREAD
+#define ELEMS_PER_THREAD 8
+#endif
+
+#define ELEMS_PER_BLOCK (THREADS_PER_WG * ELEMS_PER_THREAD)
+#define BLOCKS_PER_CHUNK (CHUNK_SIZE / ELEMS_PER_BLOCK)
+
+extern "C" __attribute__((global))
+__attribute__((amdgpu_flat_work_group_size(1, THREADS_PER_WG)))
+void fused_pad_grad_accum(
+    __hip_bfloat16* __restrict__ grad_buf
+    __FUSED_PAD_GRAD_ACCUM_PARAMS
+) {
+  const int bid = blockIdx.x;
+  const int chunk_idx = bid / BLOCKS_PER_CHUNK;
+  const int block_in_chunk = bid - chunk_idx * BLOCKS_PER_CHUNK;
+  const int tid = threadIdx.x;
+
+  const __hip_bfloat16* chunk_ptr;
+  switch (chunk_idx) {
+    __FUSED_PAD_GRAD_ACCUM_DISPATCH
+    default: chunk_ptr = (const __hip_bfloat16*)0; break;  // unreachable
+  }
+
+  // int64 for global_offset: at 32 chunks × 117M elements = 3.6B, int32 overflows → MEMVIOL.
+  const int local_offset = block_in_chunk * ELEMS_PER_BLOCK + tid * ELEMS_PER_THREAD;
+  const long long global_offset = (long long)chunk_idx * (long long)CHUNK_SIZE + (long long)local_offset;
+
+  // Vectorized 16-byte load (uint4 = 8 bf16). Requires CHUNK_SIZE % 8 == 0 and 16-byte alignment.
+  const uint4 chunk_v = *reinterpret_cast<const uint4*>(&chunk_ptr[local_offset]);
+  const uint4 grad_v  = *reinterpret_cast<const uint4*>(&grad_buf[global_offset]);
+  uint4 out_v;
+
+  const __hip_bfloat16* chunk_bf = reinterpret_cast<const __hip_bfloat16*>(&chunk_v);
+  const __hip_bfloat16* grad_bf  = reinterpret_cast<const __hip_bfloat16*>(&grad_v);
+  __hip_bfloat16*       out_bf   = reinterpret_cast<__hip_bfloat16*>(&out_v);
+
+  #pragma unroll
+  for (int i = 0; i < ELEMS_PER_THREAD; i++) {
+    out_bf[i] = (__hip_bfloat16)((float)grad_bf[i] + (float)chunk_bf[i]);
+  }
+
+  *reinterpret_cast<uint4*>(&grad_buf[global_offset]) = out_v;
+}
@@ -0,0 +1,153 @@
+from __future__ import annotations
+import functools, pathlib
+from tinygrad import Tensor, dtypes
+from tinygrad.uop.ops import UOp, Ops, KernelInfo
+from tinygrad.renderer import Estimates
+from extra.llama_kernels import FP8_MAX, NUM_WG, THREADS_PER_WG, alloc_like, alloc_local, scalar_amax, dname_of, compile_hip
+
+def _src() -> str: return (pathlib.Path(__file__).parent/"fused_rmsnorm_mul_quantize_fp8.cpp").read_text()
+def _src_bwd() -> str: return (pathlib.Path(__file__).parent/"fused_rmsnorm_mul_quantize_fp8_bwd.cpp").read_text()
+
+@functools.cache
+def _custom_fwd(fp8_out:UOp, x_normed_out:UOp, rrms_out:UOp, amax_buf:UOp,
+                x:UOp, weight:UOp, amax_state:UOp, dname:str, eps_val:float) -> UOp:
+  MBS, SEQ, HIDDEN = x.shape
+  n_elems = MBS * SEQ * HIDDEN
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
+  mem = n_elems * 2 + n_elems + MBS * SEQ * 4 + n_elems + HIDDEN * 2 + NUM_WG * 4 + 4
+  sink = UOp.sink(fp8_out.base, x_normed_out.base, rrms_out.base, amax_buf.base,
+                  x.base, weight.base, amax_state.base, threads, workgroups,
+                  arg=KernelInfo(f"fused_rmsnorm_mul_quantize_fp8_{n_elems}_h{HIDDEN}_eps{eps_val:.0e}",
+                                 estimates=Estimates(ops=6*n_elems, mem=mem)))
+  defines = [f"-DN_ELEMS={n_elems}", f"-DHIDDEN={HIDDEN}", f"-DNUM_WG={NUM_WG}", f"-DTHREADS_PER_WG={THREADS_PER_WG}",
+             f"-DEPS_LITERAL={eps_val}f"]
+  src = _src()
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=compile_hip(src, defines))))
+
+@functools.cache
+def _custom_fwd_add(fp8_out:UOp, h_out:UOp, x_normed_out:UOp, rrms_out:UOp, amax_buf:UOp,
+                    x:UOp, residual:UOp, weight:UOp, amax_state:UOp, dname:str, eps_val:float) -> UOp:
+  MBS, SEQ, HIDDEN = x.shape
+  n_elems = MBS * SEQ * HIDDEN
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
+  mem = n_elems * 2 * 4 + MBS * SEQ * 4 + HIDDEN * 2 + NUM_WG * 4 + 4
+  sink = UOp.sink(fp8_out.base, h_out.base, x_normed_out.base, rrms_out.base, amax_buf.base,
+                  x.base, residual.base, weight.base, amax_state.base, threads, workgroups,
+                  arg=KernelInfo(f"fused_add_rmsnorm_mul_quantize_fp8_{n_elems}_h{HIDDEN}_eps{eps_val:.0e}",
+                                 estimates=Estimates(ops=7*n_elems, mem=mem)))
+  defines = [f"-DN_ELEMS={n_elems}", f"-DHIDDEN={HIDDEN}", f"-DNUM_WG={NUM_WG}", f"-DTHREADS_PER_WG={THREADS_PER_WG}",
+             f"-DEPS_LITERAL={eps_val}f", f"-DHAS_RESIDUAL=1"]
+  src = _src()
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=compile_hip(src, defines))))
+
+@functools.cache
+def _custom_bwd(grad_x:UOp, grad_weight_partial:UOp,
+                grad_fp8:UOp, x_normed:UOp, rrms:UOp, weight:UOp, amax_state:UOp, dname:str) -> UOp:
+  MBS, SEQ, HIDDEN = x_normed.shape
+  n_elems = MBS * SEQ * HIDDEN
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
+  mem = n_elems * 2 * 3 + NUM_WG * HIDDEN * 4 + MBS * SEQ * 4 + HIDDEN * 2 + 4
+  sink = UOp.sink(grad_x.base, grad_weight_partial.base,
+                  grad_fp8.base, x_normed.base, rrms.base, weight.base, amax_state.base, threads, workgroups,
+                  arg=KernelInfo(f"fused_rmsnorm_mul_quantize_fp8_bwd_{n_elems}_h{HIDDEN}",
+                                 estimates=Estimates(ops=8*n_elems, mem=mem)))
+  defines = [f"-DN_ELEMS={n_elems}", f"-DHIDDEN={HIDDEN}", f"-DNUM_WG={NUM_WG}", f"-DTHREADS_PER_WG={THREADS_PER_WG}"]
+  src = _src_bwd()
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=compile_hip(src, defines))))
+
+def _bwd_common(fp8_grad_u, h_grad_u, x_u, x_normed_u, rrms_u, weight_u, amax_state_u, kernel:UOp):
+  device = x_u.device
+  MBS, SEQ, HIDDEN = x_normed_u.shape
+  axis = x_normed_u.axis if isinstance(device, tuple) else None
+  grad_x = alloc_like((MBS, SEQ, HIDDEN), dtypes.bfloat16, device, axis)
+  grad_weight_partial = alloc_local((NUM_WG, HIDDEN), dtypes.float32, device, axis)
+  grad_h_from_fp8 = None
+  grad_weight_uop = None
+  if fp8_grad_u is not None:
+    fxn = functools.partial(_custom_bwd, dname=dname_of(device))
+    grad_x_t, grad_weight_partial_t, *_ = Tensor.custom_kernel(
+      grad_x, grad_weight_partial,
+      Tensor(fp8_grad_u, device=device).cast(dtypes.bfloat16),
+      Tensor(x_normed_u.after(kernel), device=device),
+      Tensor(rrms_u.after(kernel), device=device),
+      Tensor(weight_u, device=device),
+      Tensor(amax_state_u, device=device), fxn=fxn)
+    grad_h_from_fp8 = grad_x_t
+    grad_weight_uop = grad_weight_partial_t.sum(axis=0).cast(dtypes.bfloat16).uop
+  if h_grad_u is not None:
+    h_grad_t = Tensor(h_grad_u, device=device).cast(dtypes.bfloat16)
+    grad_total = (grad_h_from_fp8 + h_grad_t) if grad_h_from_fp8 is not None else h_grad_t
+  else:
+    grad_total = grad_h_from_fp8
+  return grad_total.uop, grad_weight_uop
+
+def _fused_bwd(gradient:UOp, kernel:UOp):
+  # NOTE: fwd inputs (fp8_out, x_normed_out, rrms_out, amax_buf, x, weight, amax_state)
+  _, x_normed_u, rrms_u, _, x_u, weight_u, amax_state_u = kernel.src[1:]
+  grad_x, grad_w = _bwd_common(gradient, None, x_u, x_normed_u, rrms_u, weight_u, amax_state_u, kernel)
+  return (None, None, None, None, grad_x, grad_w, None)
+
+def _fused_add_bwd(*args, **kwargs):
+  # Two invocation modes: 1 grad => positional; >1 grads => kwarg `call=`.
+  # Outputs: (fp8_out, h_out, x_normed_out, rrms_out, amax_buf). Both fp8 and h may be consumed
+  # downstream — TUPLE order in gradient.py preserves kernel-output slot order.
+  # Don't dispatch by dtype: matmul's bwd emits fp8 grad as bf16 (no explicit cast), so
+  # dtype-detection collapses both into h_grad and silently drops the rmsnorm-bwd path.
+  if 'call' in kwargs:
+    kernel, all_grads = kwargs['call'], list(args)
+  else:
+    gradient, kernel = args
+    all_grads = [gradient]
+  fp8_grad_u = h_grad_u = None
+  if len(all_grads) >= 2:
+    fp8_grad_u, h_grad_u = all_grads[0], all_grads[1]
+  elif len(all_grads) == 1:
+    g = all_grads[0]
+    if g.dtype == dtypes.bfloat16: h_grad_u = g
+    else: fp8_grad_u = g
+  _, _, x_normed_u, rrms_u, _, x_u, _, weight_u, amax_state_u = kernel.src[1:]
+  grad_h, grad_w = _bwd_common(fp8_grad_u, h_grad_u, x_u, x_normed_u, rrms_u, weight_u, amax_state_u, kernel)
+  return (None, None, None, None, None, grad_h, grad_h, grad_w, None)
+
+def fused_rmsnorm_mul_quantize_fp8(x:Tensor, weight:Tensor, amax_state:Tensor, eps:float, fp8_dtype) -> tuple[Tensor, Tensor, Tensor, Tensor, Tensor]:
+  # NOTE: rmsnorm(x) * weight -> fp8 + amax. Returns (fp8, inv_scale, new_amax, x_normed, rrms).
+  # x_normed + rrms are saved for the rmsnorm backward (also recomputed here from x regs).
+  assert x.dtype == dtypes.bfloat16 and weight.dtype == dtypes.bfloat16
+  assert x.shape[-1] == weight.shape[-1], f"HIDDEN mismatch: x={x.shape}, weight={weight.shape}"
+  MBS, SEQ, HIDDEN = x.shape
+  axis = x.uop.axis if isinstance(x.device, tuple) else None
+  if isinstance(x.device, tuple): assert axis in (None, 0, 1), f"unsupported sharding axis={axis}"
+  fp8_out      = alloc_like((MBS, SEQ, HIDDEN), fp8_dtype,       x.device, axis)
+  x_normed_out = alloc_like((MBS, SEQ, HIDDEN), dtypes.bfloat16, x.device, axis)
+  rrms_out     = alloc_like((MBS, SEQ),         dtypes.float32,  x.device, axis)
+  amax_buf     = alloc_local((NUM_WG,),         dtypes.float32,  x.device, axis)
+  fxn = functools.partial(_custom_fwd, dname=dname_of(x.device), eps_val=eps)
+  fp8_out, x_normed_out, rrms_out, amax_buf, *_ = Tensor.custom_kernel(
+    fp8_out, x_normed_out, rrms_out, amax_buf, x, weight, amax_state, fxn=fxn, grad_fxn=_fused_bwd)
+  inv_scale = (amax_state.float() + 1e-8) / FP8_MAX
+  return fp8_out, inv_scale, scalar_amax(amax_buf), x_normed_out, rrms_out
+
+def fused_add_rmsnorm_mul_quantize_fp8(x:Tensor, residual:Tensor, weight:Tensor, amax_state:Tensor,
+                                       eps:float, fp8_dtype) -> tuple[Tensor, Tensor, Tensor, Tensor, Tensor, Tensor]:
+  # NOTE: h = x + residual; y_normed = rmsnorm(h); fp8 = quantize(y_normed * weight).
+  # Returns (fp8, inv_scale, new_amax, h, x_normed, rrms). h is also written so downstream can
+  # reuse it without recomputing x+residual — eliminates the separate residual-add kernel.
+  assert x.dtype == dtypes.bfloat16 and residual.dtype == dtypes.bfloat16 and weight.dtype == dtypes.bfloat16
+  assert x.shape == residual.shape
+  MBS, SEQ, HIDDEN = x.shape
+  axis = x.uop.axis if isinstance(x.device, tuple) else None
+  if isinstance(x.device, tuple): assert axis in (None, 0, 1), f"unsupported sharding axis={axis}"
+  fp8_out      = alloc_like((MBS, SEQ, HIDDEN), fp8_dtype,       x.device, axis)
+  h_out        = alloc_like((MBS, SEQ, HIDDEN), dtypes.bfloat16, x.device, axis)
+  x_normed_out = alloc_like((MBS, SEQ, HIDDEN), dtypes.bfloat16, x.device, axis)
+  rrms_out     = alloc_like((MBS, SEQ),         dtypes.float32,  x.device, axis)
+  amax_buf     = alloc_local((NUM_WG,),         dtypes.float32,  x.device, axis)
+  fxn = functools.partial(_custom_fwd_add, dname=dname_of(x.device), eps_val=eps)
+  fp8_out, h_out, x_normed_out, rrms_out, amax_buf, *_ = Tensor.custom_kernel(
+    fp8_out, h_out, x_normed_out, rrms_out, amax_buf, x, residual, weight, amax_state,
+    fxn=fxn, grad_fxn=_fused_add_bwd)
+  inv_scale = (amax_state.float() + 1e-8) / FP8_MAX
+  return fp8_out, inv_scale, scalar_amax(amax_buf), h_out, x_normed_out, rrms_out
@@ -0,0 +1,155 @@
+#include <hip/hip_runtime.h>
+#include <hip/hip_bf16.h>
+#include <hip/hip_fp8.h>
+
+// Fuses the full pre-matmul preparation for a layer into a single HBM pass:
+//   y = rmsnorm(x) * weight            (reduce-mean-square + rsqrt + per-elem mul)
+//   fp8 = fp8_sat(y * (FP8_MAX / amax_state))
+// Also writes:
+//   rrms[row]        — saved for the rmsnorm backward
+//   amax_buf[wg]     — per-WG |y| partials, reduced later to update amax_state
+//
+// Layout: one WG per row, ROWS_PER_WG rows per WG via grid-stride (ROWS = N_ELEMS / HIDDEN).
+// Each thread handles HIDDEN / THREADS_PER_WG elements per row.
+
+#ifndef N_ELEMS
+#define N_ELEMS 67108864
+#endif
+#ifndef HIDDEN
+#define HIDDEN 4096
+#endif
+#ifndef NUM_WG
+#define NUM_WG 1024
+#endif
+#ifndef THREADS_PER_WG
+#define THREADS_PER_WG 256
+#endif
+#ifndef EPS_LITERAL
+#define EPS_LITERAL 1e-5f
+#endif
+#ifndef HAS_RESIDUAL
+#define HAS_RESIDUAL 0
+#endif
+
+constexpr int VEC = 8;
+constexpr float FP8_MAX = 448.0f;
+
+static_assert(N_ELEMS % HIDDEN == 0, "N_ELEMS must be a multiple of HIDDEN");
+static_assert(HIDDEN % (THREADS_PER_WG * VEC) == 0, "HIDDEN must be divisible by THREADS_PER_WG*VEC");
+
+constexpr int ROWS = N_ELEMS / HIDDEN;
+constexpr int ELEMS_PER_THREAD = HIDDEN / THREADS_PER_WG;  // each thread sees this many elems per row
+constexpr int VECS_PER_THREAD = ELEMS_PER_THREAD / VEC;    // number of 8-wide vec loads
+
+#if HAS_RESIDUAL
+extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
+fused_add_rmsnorm_mul_quantize_fp8(
+    __hip_fp8_storage_t*  __restrict__ fp8_out,         // fp8, ROWS*HIDDEN
+    __hip_bfloat16*       __restrict__ h_out,           // bf16, ROWS*HIDDEN — x + residual (saved for downstream)
+    __hip_bfloat16*       __restrict__ x_normed_out,    // bf16, ROWS*HIDDEN
+    float*                __restrict__ rrms_out,        // fp32, ROWS
+    float*                __restrict__ amax_buf,        // fp32, NUM_WG
+    const __hip_bfloat16* __restrict__ x,               // bf16, ROWS*HIDDEN
+    const __hip_bfloat16* __restrict__ residual,        // bf16, ROWS*HIDDEN — added into x before rmsnorm
+    const __hip_bfloat16* __restrict__ weight,          // bf16, HIDDEN
+    const float*          __restrict__ amax_state)      // fp32 scalar
+{
+#else
+extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
+fused_rmsnorm_mul_quantize_fp8(
+    __hip_fp8_storage_t*  __restrict__ fp8_out,         // fp8, ROWS*HIDDEN
+    __hip_bfloat16*       __restrict__ x_normed_out,    // bf16, ROWS*HIDDEN (saved for rmsnorm bwd)
+    float*                __restrict__ rrms_out,        // fp32, ROWS (fp32 to match rmsnorm_bwd.cpp expectation)
+    float*                __restrict__ amax_buf,        // fp32, NUM_WG per-WG partials
+    const __hip_bfloat16* __restrict__ x,               // bf16, ROWS*HIDDEN
+    const __hip_bfloat16* __restrict__ weight,          // bf16, HIDDEN (per-hidden scale)
+    const float*          __restrict__ amax_state)      // fp32 scalar
+{
+#endif
+  __shared__ float sdata[THREADS_PER_WG];
+
+  const int tid = threadIdx.x;
+  const int wg  = blockIdx.x;
+
+  const float scale = FP8_MAX / (static_cast<float>(*amax_state) + 1e-8f);
+  const float inv_hidden = 1.0f / static_cast<float>(HIDDEN);
+  float local_max = 0.0f;
+
+  // Grid-stride over rows. Each WG processes rows (wg, wg+NUM_WG, wg+2*NUM_WG, ...).
+  for (int row = wg; row < ROWS; row += NUM_WG) {
+    const int row_off = row * HIDDEN;
+
+    // Load row (+ residual if present) into registers.
+    float regs[ELEMS_PER_THREAD];
+    float sum_sq = 0.0f;
+    #pragma unroll
+    for (int v = 0; v < VECS_PER_THREAD; v++) {
+      const int h_base = tid * VEC + v * THREADS_PER_WG * VEC;
+      float4 raw = *reinterpret_cast<const float4*>(&x[row_off + h_base]);
+      const __hip_bfloat16 *xi = reinterpret_cast<const __hip_bfloat16*>(&raw);
+#if HAS_RESIDUAL
+      float4 res_raw = *reinterpret_cast<const float4*>(&residual[row_off + h_base]);
+      const __hip_bfloat16 *ri = reinterpret_cast<const __hip_bfloat16*>(&res_raw);
+      __hip_bfloat16 h_buf[VEC];
+#endif
+      #pragma unroll
+      for (int i = 0; i < VEC; i++) {
+#if HAS_RESIDUAL
+        const float f = static_cast<float>(xi[i]) + static_cast<float>(ri[i]);
+        h_buf[i] = static_cast<__hip_bfloat16>(f);
+#else
+        const float f = static_cast<float>(xi[i]);
+#endif
+        regs[v * VEC + i] = f;
+        sum_sq += f * f;
+      }
+#if HAS_RESIDUAL
+      *reinterpret_cast<float4*>(&h_out[row_off + h_base]) = *reinterpret_cast<float4*>(h_buf);
+#endif
+    }
+
+    // LDS tree-reduce sum_sq across the WG.
+    sdata[tid] = sum_sq;
+    __syncthreads();
+    for (int s = THREADS_PER_WG / 2; s > 0; s >>= 1) {
+      if (tid < s) sdata[tid] = sdata[tid] + sdata[tid + s];
+      __syncthreads();
+    }
+    const float mean_sq = sdata[0] * inv_hidden;
+    const float rrms = 1.0f / sqrtf(mean_sq + EPS_LITERAL);
+
+    if (tid == 0) rrms_out[row] = rrms;
+
+    // Normalize, multiply by weight, quantize. Also write x_normed (for rmsnorm bwd).
+    #pragma unroll
+    for (int v = 0; v < VECS_PER_THREAD; v++) {
+      const int h_base = tid * VEC + v * THREADS_PER_WG * VEC;
+      float4 w_raw = *reinterpret_cast<const float4*>(&weight[h_base]);
+      const __hip_bfloat16 *wi = reinterpret_cast<const __hip_bfloat16*>(&w_raw);
+
+      __hip_fp8_storage_t out[VEC];
+      __hip_bfloat16 xn[VEC];
+      #pragma unroll
+      for (int i = 0; i < VEC; i++) {
+        const float x_normed = regs[v * VEC + i] * rrms;
+        xn[i] = static_cast<__hip_bfloat16>(x_normed);
+        const float y = x_normed * static_cast<float>(wi[i]);
+        local_max = fmaxf(local_max, fabsf(y));
+        const float scaled = fmaxf(-FP8_MAX, fminf(FP8_MAX, y * scale));
+        out[i] = __hip_cvt_float_to_fp8(scaled, __HIP_SATFINITE, __HIP_E4M3);
+      }
+      *reinterpret_cast<uint64_t*>(&fp8_out[row_off + h_base]) = *reinterpret_cast<uint64_t*>(out);
+      *reinterpret_cast<float4*>(&x_normed_out[row_off + h_base]) = *reinterpret_cast<float4*>(xn);
+    }
+    __syncthreads();  // before next row's sum_sq reduce reuses sdata
+  }
+
+  // Final per-WG amax reduce.
+  sdata[tid] = local_max;
+  __syncthreads();
+  for (int s = THREADS_PER_WG / 2; s > 0; s >>= 1) {
+    if (tid < s) sdata[tid] = fmaxf(sdata[tid], sdata[tid + s]);
+    __syncthreads();
+  }
+  if (tid == 0) amax_buf[wg] = sdata[0];
+}
@@ -0,0 +1,147 @@
+#include <hip/hip_runtime.h>
+#include <hip/hip_bf16.h>
+
+// Full backward for fused_rmsnorm_mul_quantize_fp8.cpp. One HBM pass per row produces:
+//   grad_x (bf16)              — gradient w.r.t. pre-rmsnorm x
+//   grad_weight_partial (fp32) — per-WG partial of the weight gradient, reduced later
+//
+// Input (all read):
+//   grad_fp8 (bf16)     — upstream grad w.r.t. fp8_out (bf16-typed gradient value)
+//   x_normed (bf16)     — saved from the fwd kernel, shape (ROWS, HIDDEN)
+//   rrms (fp32)         — saved rrms per row
+//   weight (bf16)       — per-HIDDEN rmsnorm weight
+//   amax_state (bf16)   — delayed amax used to compute the fp8 scale in fwd
+//
+// Chain: y = x_normed * weight; fp8 = sat(y * scale). Through STE: grad_y = grad_fp8 * scale.
+//        grad_x_normed = grad_y * weight.
+//        grad_weight   = sum_rows(grad_y * x_normed).
+//        grad_x        = rrms * (grad_x_normed - x_normed * mean(grad_x_normed * x_normed, last_dim)).
+
+#ifndef N_ELEMS
+#define N_ELEMS 67108864
+#endif
+#ifndef HIDDEN
+#define HIDDEN 4096
+#endif
+#ifndef NUM_WG
+#define NUM_WG 1024
+#endif
+#ifndef THREADS_PER_WG
+#define THREADS_PER_WG 256
+#endif
+
+constexpr int VEC = 8;
+constexpr float FP8_MAX = 448.0f;
+
+static_assert(N_ELEMS % HIDDEN == 0, "N_ELEMS must be a multiple of HIDDEN");
+static_assert(HIDDEN % (THREADS_PER_WG * VEC) == 0, "HIDDEN must be divisible by THREADS_PER_WG*VEC");
+
+constexpr int ROWS = N_ELEMS / HIDDEN;
+constexpr int ELEMS_PER_THREAD = HIDDEN / THREADS_PER_WG;
+constexpr int VECS_PER_THREAD = ELEMS_PER_THREAD / VEC;
+
+extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
+fused_rmsnorm_mul_quantize_fp8_bwd(
+    __hip_bfloat16*       __restrict__ grad_x,              // out: bf16, ROWS*HIDDEN
+    float*                __restrict__ grad_weight_partial, // out: fp32, NUM_WG*HIDDEN
+    const __hip_bfloat16* __restrict__ grad_fp8,            // in:  bf16, ROWS*HIDDEN (grad of fp8_out)
+    const __hip_bfloat16* __restrict__ x_normed,            // in:  bf16, ROWS*HIDDEN
+    const float*          __restrict__ rrms,                // in:  fp32, ROWS
+    const __hip_bfloat16* __restrict__ weight,              // in:  bf16, HIDDEN
+    const float*          __restrict__ amax_state)          // in:  fp32 scalar
+{
+  __shared__ float sdata[THREADS_PER_WG];
+
+  const int tid = threadIdx.x;
+  const int wg  = blockIdx.x;
+
+  const float scale = FP8_MAX / (static_cast<float>(*amax_state) + 1e-8f);
+  const float inv_hidden = 1.0f / static_cast<float>(HIDDEN);
+
+  // Per-thread accumulator for grad_weight (across all rows this WG touches).
+  float gw_accum[ELEMS_PER_THREAD];
+  #pragma unroll
+  for (int i = 0; i < ELEMS_PER_THREAD; i++) gw_accum[i] = 0.0f;
+
+  // Preload weight into registers (same across rows). Use ELEMS_PER_THREAD entries.
+  float w_regs[ELEMS_PER_THREAD];
+  #pragma unroll
+  for (int v = 0; v < VECS_PER_THREAD; v++) {
+    const int h_base = tid * VEC + v * THREADS_PER_WG * VEC;
+    float4 w_raw = *reinterpret_cast<const float4*>(&weight[h_base]);
+    const __hip_bfloat16 *wi = reinterpret_cast<const __hip_bfloat16*>(&w_raw);
+    #pragma unroll
+    for (int i = 0; i < VEC; i++) w_regs[v * VEC + i] = static_cast<float>(wi[i]);
+  }
+
+  for (int row = wg; row < ROWS; row += NUM_WG) {
+    const int row_off = row * HIDDEN;
+    const float rrms_v = rrms[row];
+
+    // Load grad_fp8 and x_normed rows into registers, compute grad_y and grad_x_normed.
+    float g_y_regs[ELEMS_PER_THREAD];
+    float xn_regs[ELEMS_PER_THREAD];
+    float g_xn_regs[ELEMS_PER_THREAD];  // grad_x_normed
+    float local_dot = 0.0f;  // sum(grad_x_normed * x_normed) for mean
+
+    #pragma unroll
+    for (int v = 0; v < VECS_PER_THREAD; v++) {
+      const int h_base = tid * VEC + v * THREADS_PER_WG * VEC;
+      float4 g_raw  = *reinterpret_cast<const float4*>(&grad_fp8[row_off + h_base]);
+      float4 xn_raw = *reinterpret_cast<const float4*>(&x_normed[row_off + h_base]);
+      const __hip_bfloat16 *gi  = reinterpret_cast<const __hip_bfloat16*>(&g_raw);
+      const __hip_bfloat16 *xni = reinterpret_cast<const __hip_bfloat16*>(&xn_raw);
+
+      #pragma unroll
+      for (int i = 0; i < VEC; i++) {
+        const int idx = v * VEC + i;
+        const float g_y = static_cast<float>(gi[i]) * scale;
+        const float xn  = static_cast<float>(xni[i]);
+        g_y_regs[idx] = g_y;
+        xn_regs[idx]  = xn;
+        g_xn_regs[idx] = g_y * w_regs[idx];           // grad_x_normed = grad_y * weight
+        gw_accum[idx] += g_y * xn;                    // grad_weight contrib
+        local_dot += g_xn_regs[idx] * xn;             // for mean
+      }
+    }
+
+    // LDS reduce local_dot to sdata[0].
+    sdata[tid] = local_dot;
+    __syncthreads();
+    for (int s = THREADS_PER_WG / 2; s > 0; s >>= 1) {
+      if (tid < s) sdata[tid] = sdata[tid] + sdata[tid + s];
+      __syncthreads();
+    }
+    const float mean_term = sdata[0] * inv_hidden;
+
+    // Compute grad_x = rrms * (grad_x_normed - x_normed * mean_term) and write.
+    #pragma unroll
+    for (int v = 0; v < VECS_PER_THREAD; v++) {
+      const int h_base = tid * VEC + v * THREADS_PER_WG * VEC;
+      __hip_bfloat16 out[VEC];
+      #pragma unroll
+      for (int i = 0; i < VEC; i++) {
+        const int idx = v * VEC + i;
+        const float dx = rrms_v * (g_xn_regs[idx] - xn_regs[idx] * mean_term);
+        out[i] = static_cast<__hip_bfloat16>(dx);
+      }
+      *reinterpret_cast<float4*>(&grad_x[row_off + h_base]) = *reinterpret_cast<float4*>(out);
+    }
+    __syncthreads();
+  }
+
+  // Write this WG's grad_weight partial to HBM (fp32, NUM_WG x HIDDEN layout).
+  const int gw_row_off = wg * HIDDEN;
+  #pragma unroll
+  for (int v = 0; v < VECS_PER_THREAD; v++) {
+    const int h_base = tid * VEC + v * THREADS_PER_WG * VEC;
+    // Write 8 fp32 values with two float4 stores.
+    float4 out_lo, out_hi;
+    out_lo.x = gw_accum[v * VEC + 0]; out_lo.y = gw_accum[v * VEC + 1];
+    out_lo.z = gw_accum[v * VEC + 2]; out_lo.w = gw_accum[v * VEC + 3];
+    out_hi.x = gw_accum[v * VEC + 4]; out_hi.y = gw_accum[v * VEC + 5];
+    out_hi.z = gw_accum[v * VEC + 6]; out_hi.w = gw_accum[v * VEC + 7];
+    *reinterpret_cast<float4*>(&grad_weight_partial[gw_row_off + h_base + 0]) = out_lo;
+    *reinterpret_cast<float4*>(&grad_weight_partial[gw_row_off + h_base + 4]) = out_hi;
+  }
+}
@@ -0,0 +1,67 @@
+from __future__ import annotations
+import functools, pathlib
+from tinygrad import Tensor, dtypes
+from tinygrad.uop.ops import UOp, Ops, KernelInfo
+from tinygrad.renderer import Estimates
+from extra.llama_kernels import FP8_MAX, NUM_WG, THREADS_PER_WG, alloc_like, alloc_local, scalar_amax, dname_of, compile_hip
+
+@functools.cache
+def _custom_quantize_fp8_with_amax(fp8_out:UOp, amax_partial:UOp, x:UOp, amax_state:UOp, dname:str) -> UOp:
+  n_elems = 1
+  for d in x.shape: n_elems *= d
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
+  mem = n_elems * 2 + n_elems + 4 + NUM_WG * 4
+  sink = UOp.sink(fp8_out.base, amax_partial.base, x.base, amax_state.base, threads, workgroups,
+                  arg=KernelInfo(f"quantize_fp8_with_amax_{n_elems}", estimates=Estimates(ops=3*n_elems, mem=mem)))
+  src = (pathlib.Path(__file__).parent/"quantize_fp8_with_amax.cpp").read_text()
+  defines = [f"-DN_ELEMS={n_elems}", f"-DNUM_WG={NUM_WG}", f"-DTHREADS_PER_WG={THREADS_PER_WG}"]
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=compile_hip(src, defines))))
+
+@functools.cache
+def _custom_quantize_fp8_scalar(fp8_out:UOp, x:UOp, amax_state:UOp, dname:str) -> UOp:
+  n_elems = 1
+  for d in x.shape: n_elems *= d
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
+  mem = n_elems * 2 + n_elems
+  sink = UOp.sink(fp8_out.base, x.base, amax_state.base, threads, workgroups,
+                  arg=KernelInfo(f"quantize_fp8_scalar_{n_elems}", estimates=Estimates(ops=2*n_elems, mem=mem)))
+  src = (pathlib.Path(__file__).parent/"quantize_fp8_scalar.cpp").read_text()
+  defines = [f"-DN_ELEMS={n_elems}", f"-DNUM_WG={NUM_WG}", f"-DTHREADS_PER_WG={THREADS_PER_WG}"]
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=compile_hip(src, defines))))
+
+def _quantize_fp8_delayed_bwd(gradient:UOp, kernel:UOp):
+  # NOTE: STE-equivalent backward — grad_x = grad_fp8 * scale, scale = FP8_MAX / amax_state.
+  # `gradient` is bf16 grad w.r.t. fp8 output (asm_gemm bwd already applied x_scale).
+  _, _, x, amax_state = kernel.src[1:]
+  device = x.device
+  scale = FP8_MAX / (Tensor(amax_state, device=device).float() + 1e-8)
+  grad_x = (Tensor(gradient, device=device).float() * scale).cast(dtypes.bfloat16)
+  return (None, None, grad_x.uop, None)
+
+def quantize_fp8_delayed(x:Tensor, amax_state:Tensor, fp8_dtype=dtypes.fp8e4m3) -> tuple[Tensor, Tensor, Tensor, UOp]:
+  # NOTE: one-pass bf16 -> fp8 quantize with delayed scaling. Returns (fp8, inv_scale, new_amax, store_effect).
+  # Fused kernel reads x once and writes fp8 + per-WG |x| partials (then a small reduce produces scalar new_amax).
+  # store_effect writes new_amax into amax_state's buffer — the caller must thread it into a realized
+  # output via `.after(store_effect)`. Calling `amax_state.assign(new_amax)` inside a grad_fxn does
+  # NOT work because .assign mutates only the temp Tensor's .uop, not the original layer-owned buffer.
+  assert x.dtype == dtypes.bfloat16, f"expected bf16, got {x.dtype}"
+  axis = x.uop.axis if isinstance(x.device, tuple) else None
+  fp8_out      = alloc_like(x.shape,  fp8_dtype,      x.device, axis)
+  amax_partial = alloc_local((NUM_WG,), dtypes.float32, x.device, axis)
+  fxn = functools.partial(_custom_quantize_fp8_with_amax, dname=dname_of(x.device))
+  fp8_out, amax_partial, *_ = Tensor.custom_kernel(fp8_out, amax_partial, x, amax_state,
+                                                    fxn=fxn, grad_fxn=_quantize_fp8_delayed_bwd)
+  new_amax = scalar_amax(amax_partial)
+  inv_scale = (amax_state.float() + 1e-8) / FP8_MAX
+  store_effect = amax_state.uop.store(new_amax.uop)
+  return fp8_out, inv_scale, new_amax, store_effect
+
+def quantize_fp8_scalar(x:Tensor, amax_state:Tensor, fp8_dtype=dtypes.fp8e4m3) -> Tensor:
+  # NOTE: pure one-pass bf16 -> fp8 quantize with delayed scalar scale. No amax computation.
+  axis = x.uop.axis if isinstance(x.device, tuple) else None
+  fp8_out = alloc_like(x.shape, fp8_dtype, x.device, axis)
+  fxn = functools.partial(_custom_quantize_fp8_scalar, dname=dname_of(x.device))
+  fp8_out, *_ = Tensor.custom_kernel(fp8_out, x, amax_state, fxn=fxn)
+  return fp8_out
@@ -0,0 +1,48 @@
+#include <hip/hip_runtime.h>
+#include <hip/hip_bf16.h>
+#include <hip/hip_fp8.h>
+
+// Pure one-pass bf16 -> fp8 quantize with delayed scalar scale. No amax computation.
+
+#ifndef N_ELEMS
+#define N_ELEMS 67108864
+#endif
+#ifndef NUM_WG
+#define NUM_WG 1024
+#endif
+#ifndef THREADS_PER_WG
+#define THREADS_PER_WG 256
+#endif
+
+constexpr int VEC = 8;
+constexpr float FP8_MAX = 448.0f;
+
+static_assert(N_ELEMS % VEC == 0, "N_ELEMS must be divisible by VEC");
+
+extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
+quantize_fp8_scalar(
+    __hip_fp8_storage_t*  __restrict__ fp8_out,         // fp8, N_ELEMS
+    const __hip_bfloat16* __restrict__ x,               // bf16, N_ELEMS
+    const float*          __restrict__ amax_state)      // fp32 scalar (delayed)
+{
+  const int tid = threadIdx.x;
+  const int wg  = blockIdx.x;
+  const int gid = wg * THREADS_PER_WG + tid;
+  const int stride_elems = NUM_WG * THREADS_PER_WG * VEC;
+
+  const float scale = FP8_MAX / (static_cast<float>(*amax_state) + 1e-8f);
+
+  for (int base = gid * VEC; base < N_ELEMS; base += stride_elems) {
+    float4 x_raw = *reinterpret_cast<const float4*>(&x[base]);
+    const __hip_bfloat16 *xi = reinterpret_cast<const __hip_bfloat16*>(&x_raw);
+
+    __hip_fp8_storage_t out[VEC];
+    #pragma unroll
+    for (int i = 0; i < VEC; i++) {
+      const float v = static_cast<float>(xi[i]);
+      const float scaled = fmaxf(-FP8_MAX, fminf(FP8_MAX, v * scale));
+      out[i] = __hip_cvt_float_to_fp8(scaled, __HIP_SATFINITE, __HIP_E4M3);
+    }
+    *reinterpret_cast<uint64_t*>(&fp8_out[base]) = *reinterpret_cast<uint64_t*>(out);
+  }
+}
@@ -0,0 +1,63 @@
+#include <hip/hip_runtime.h>
+#include <hip/hip_bf16.h>
+#include <hip/hip_fp8.h>
+
+// One-pass bf16 -> fp8 quantize using a scalar delayed amax state,
+// AND simultaneously computes per-WG |x| max partials for the next step's amax state.
+// Saves one full HBM pass over the grad tensor vs. doing quantize + separate abs().max().
+
+#ifndef N_ELEMS
+#define N_ELEMS 67108864
+#endif
+#ifndef NUM_WG
+#define NUM_WG 1024
+#endif
+#ifndef THREADS_PER_WG
+#define THREADS_PER_WG 256
+#endif
+
+constexpr int VEC = 8;
+constexpr float FP8_MAX = 448.0f;
+
+static_assert(N_ELEMS % VEC == 0, "N_ELEMS must be divisible by VEC");
+
+extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
+quantize_fp8_with_amax(
+    __hip_fp8_storage_t*  __restrict__ fp8_out,       // out: fp8, N_ELEMS
+    float*                __restrict__ amax_partial,  // out: fp32, NUM_WG per-WG partials
+    const __hip_bfloat16* __restrict__ x,             // in:  bf16, N_ELEMS
+    const float*          __restrict__ amax_state)    // in:  fp32 scalar (delayed)
+{
+  __shared__ float sdata[THREADS_PER_WG];
+
+  const int tid = threadIdx.x;
+  const int wg  = blockIdx.x;
+  const int gid = wg * THREADS_PER_WG + tid;
+  const int stride_elems = NUM_WG * THREADS_PER_WG * VEC;
+
+  const float scale = FP8_MAX / (static_cast<float>(*amax_state) + 1e-8f);
+  float local_max = 0.0f;
+
+  for (int base = gid * VEC; base < N_ELEMS; base += stride_elems) {
+    float4 x_raw = *reinterpret_cast<const float4*>(&x[base]);
+    const __hip_bfloat16 *xi = reinterpret_cast<const __hip_bfloat16*>(&x_raw);
+
+    __hip_fp8_storage_t out[VEC];
+    #pragma unroll
+    for (int i = 0; i < VEC; i++) {
+      const float v = static_cast<float>(xi[i]);
+      local_max = fmaxf(local_max, fabsf(v));
+      const float scaled = fmaxf(-FP8_MAX, fminf(FP8_MAX, v * scale));
+      out[i] = __hip_cvt_float_to_fp8(scaled, __HIP_SATFINITE, __HIP_E4M3);
+    }
+    *reinterpret_cast<uint64_t*>(&fp8_out[base]) = *reinterpret_cast<uint64_t*>(out);
+  }
+
+  sdata[tid] = local_max;
+  __syncthreads();
+  for (int s = THREADS_PER_WG / 2; s > 0; s >>= 1) {
+    if (tid < s) sdata[tid] = fmaxf(sdata[tid], sdata[tid + s]);
+    __syncthreads();
+  }
+  if (tid == 0) amax_partial[wg] = sdata[0];
+}
@@ -0,0 +1,24 @@
+from __future__ import annotations
+import functools
+from tinygrad import Tensor
+from tinygrad.uop.ops import UOp
+
+def rmsnorm_fwd(x_in:Tensor, eps:float) -> tuple[Tensor, Tensor]:
+  x = x_in.float()
+  rrms = (x.square().mean(-1, keepdim=True) + eps).rsqrt()
+  return (x * rrms).cast(x_in.dtype), rrms
+
+@functools.cache
+def _rmsnorm_fwd_fxn(x_in_p, eps, device):
+  return rmsnorm_fwd(Tensor(x_in_p, device=device), eps)
+
+def _rmsnorm_bwd(grad:UOp, call:UOp) -> tuple:
+  x_normed = Tensor(call.gettuple(0)).float()
+  do_float = Tensor(grad).float()
+  d_x = Tensor(call.gettuple(1)) * (do_float - x_normed * (do_float * x_normed).mean(-1, keepdim=True))
+  return (d_x.cast(call.src[1].dtype).uop,)
+
+def rmsnorm(x_in:Tensor, eps:float) -> tuple[Tensor, Tensor]:
+  fxn = _rmsnorm_fwd_fxn(x_in.as_param(0).uop, eps, x_in.device)
+  call = UOp.maketuple(fxn[0].uop, fxn[1].uop).call(x_in.uop, grad_fxn=_rmsnorm_bwd)
+  return Tensor(call.gettuple(0)), Tensor(call.gettuple(1))
@@ -0,0 +1,68 @@
+#!/usr/bin/env python3
+import subprocess, json, sys, os
+
+REMOTE_HOST = os.getenv("REMOTE_HOST", "192.168.52.154")
+LOCAL_PCI   = os.getenv("MLX_PCI", "0000:41:00.0")
+REMOTE_PCI  = os.getenv("REMOTE_PCI", "0000:41:00.0")
+LOCAL_IP    = os.getenv("LOCAL_IP", "10.0.0.1")
+REMOTE_IP   = os.getenv("REMOTE_IP", "10.0.0.2")
+SSH         = ["ssh", "-o", "StrictHostKeyChecking=no", REMOTE_HOST]
+TINYGRAD    = os.path.dirname(os.path.abspath(__file__)) + "/../.."
+
+print("syncing code to remote")
+subprocess.run(["rsync", "-az", "--exclude=.git", "--exclude=__pycache__", "--exclude=*.pyc",
+                TINYGRAD + "/", f"{REMOTE_HOST}:~/tinygrad/"], check=True)
+
+print("booting remote")
+remote = subprocess.Popen(
+  SSH + [f"cd ~/tinygrad && sudo PYTHONPATH=. MLX_DEBUG=1 MLX_PCI={REMOTE_PCI} MLX_IP={REMOTE_IP} python3 extra/mlx_driver/mlxdev.py --server"],
+  stdin=subprocess.PIPE, stdout=subprocess.PIPE, stderr=sys.stderr, text=True)
+
+remote_info = None
+for line in iter(remote.stdout.readline, ''):
+  print(f"  [remote] {line}", end='')
+  try: remote_info = json.loads(line.strip()); break
+  except json.JSONDecodeError: pass
+assert remote_info, "failed to get remote connection info"
+
+print("booting local")
+sys.path.insert(0, os.path.join(os.path.dirname(os.path.abspath(__file__)), "../.."))
+from extra.mlx_driver.mlxdev import MLXDev, MLXQP
+from tinygrad.runtime.support.system import PCIDevice
+
+local_dev = MLXDev(PCIDevice("mlx5", LOCAL_PCI), ip=LOCAL_IP)
+local_qp = MLXQP(local_dev)
+local_info = {"qpn": local_qp.qpn, "mac": local_dev.mac.to_bytes(6,'big').hex(), "gid": local_dev.local_gid.hex()}
+
+remote.stdin.write(json.dumps(local_info) + "\n")
+remote.stdin.flush()
+for line in iter(remote.stdout.readline, ''):
+  print(f"  [remote] {line}", end='')
+  if "connected" in line: break
+
+local_qp.connect(remote_info["qpn"], int(remote_info["mac"], 16), int(remote_info["gid"], 16))
+print("both QPs in RTS")
+
+remote_target = None
+for line in iter(remote.stdout.readline, ''):
+  print(f"  [remote] {line}", end='')
+  try: remote_target = json.loads(line.strip()); break
+  except json.JSONDecodeError: pass
+assert remote_target
+
+test_msg = b"Test message, rdma works!"
+src_mem, src_paddrs = local_dev.pci_dev.alloc_sysmem(0x1000)
+for i, b in enumerate(test_msg): src_mem[i] = b
+
+print(f"RDMA WRITE {len(test_msg)}B to remote phys 0x{remote_target['target_addr']:x}")
+local_qp.rdma_write(remote_target["target_addr"], remote_target["rkey"], src_paddrs[0], local_dev.mkey, len(test_msg))
+
+remote.stdin.write("done\n")
+remote.stdin.flush()
+for line in iter(remote.stdout.readline, ''):
+  print(f"  [remote] {line}", end='')
+  if "AS TEXT" in line: break
+
+remote.stdin.close()
+remote.wait()
+print("RDMA WRITE test complete")
@@ -0,0 +1,99 @@
+#!/usr/bin/env python3
+# GMMU=0 MLX_PCI=0000:41:00.0 PYTHONPATH=. python3 extra/mlx_driver/loopback.py
+import struct
+from tinygrad.helpers import getenv, round_up
+from tinygrad.device import Device, BufferSpec
+from tinygrad.runtime.support.system import PCIDevice
+from tinygrad.runtime.support.memory import AddrSpace
+from tinygrad.runtime.ops_amd import AMDComputeQueue
+from tinygrad.helpers import to_be32, to_be64
+from extra.mlx_driver.mlxdev import MLXDev, MLXQP
+
+BUF_SIZE = 0x1000
+MLX_PCI = getenv("MLX_PCI", "0000:41:00.0")
+MLX_IP = getenv("MLX_IP", "10.0.0.1")
+
+def map_phys_to_gpu(gpu, paddr, size):
+  size = round_up(size, 0x1000)
+  va = gpu.iface.dev_impl.mm.alloc_vaddr(size, align=0x1000)
+  gpu.iface.dev_impl.mm.map_range(va, size, [(paddr, size)], aspace=AddrSpace.SYS, snooped=True, uncached=True)
+  return va
+
+print("[init] AMD GPU...")
+gpu = Device["AMD"]
+
+print(f"[init] MLX5 at {MLX_PCI}")
+dev = MLXDev(PCIDevice("mlx5", MLX_PCI), ip=MLX_IP)
+qp = MLXQP(dev)
+
+print(f"[init] loopback connect QP 0x{qp.qp_info['qpn']:x}")
+qp.connect(qp.qp_info['qpn'], dev.mac, int.from_bytes(dev.local_gid, 'big'))
+
+# allocate src/dst via AMD GPU allocator
+buf_src = gpu.allocator.alloc(BUF_SIZE, BufferSpec(nolru=True))
+buf_dst = gpu.allocator.alloc(BUF_SIZE, BufferSpec(nolru=True))
+
+bar_base = gpu.iface.pci_dev.bar_info(gpu.iface.vram_bar)[0]
+src_paddr = buf_src.meta.mapping.paddrs[0][0] + bar_base
+dst_paddr = buf_dst.meta.mapping.paddrs[0][0] + bar_base
+print(f"src paddr=0x{src_paddr:x} dst paddr=0x{dst_paddr:x}")
+
+# fill src, zero dst
+test_msg = b"Hello from loopback send/recv!"
+gpu.allocator._copyin(buf_src, memoryview(bytearray(test_msg.ljust(BUF_SIZE, b'\x00'))))
+gpu.allocator._copyin(buf_dst, memoryview(bytearray(BUF_SIZE)))
+gpu.synchronize()
+
+# post recv WQE on RQ from CPU (scatter entry: byte_count, lkey, addr)
+rq_mask = (1 << 4) - 1  # log_rq_size=4
+rq_wqe = qp.qp_buf.view((qp.rq_head & rq_mask) * 16, 16)
+rq_wqe[:] = struct.pack('>IIQ', len(test_msg), dev.mkey, dst_paddr)
+qp.rq_head += 1
+# ring recv doorbell from CPU (DBR offset 0 = recv counter)
+dev.dbr[qp.qp_dbr // 4] = to_be32(qp.rq_head)
+
+# build send WQE in SQ from CPU (opcode 0x0a = SEND, ds_count=2)
+sq_head = qp.sq_head
+sq_mask = (1 << qp.log_sq_size) - 1
+wqe = qp.qp_buf.view(qp.sq_offset + (sq_head & sq_mask) * 64, 64)
+wqe[:] = bytes(64)
+wqe[0:8] = struct.pack('>II', (sq_head << 8) | 0x0a, (qp.qp_info['qpn'] << 8) | 2)
+wqe[11] = 0x08  # CE: signal completion
+wqe[16:32] = struct.pack('>IIQ', len(test_msg), dev.mkey, src_paddr)
+qp.sq_head += 1
+doorbell_val = to_be64(int.from_bytes(bytes(wqe[0:8]), 'big'))
+
+# map MLX5 UAR and DBR into GPU VA
+uar_paddr = dev.pci_dev.bar_info(0)[0] + dev.uar * 0x1000
+uar_gpu_va = map_phys_to_gpu(gpu, uar_paddr, 0x1000)
+dbr_gpu_va = map_phys_to_gpu(gpu, dev.dbr_paddrs[0], 0x1000)
+print(f"UAR gpu_va=0x{uar_gpu_va:x} DBR gpu_va=0x{dbr_gpu_va:x}")
+
+# GPU rings send doorbell via compute queue release_mem
+q = AMDComputeQueue(gpu)
+q.wait(gpu.timeline_signal, gpu.timeline_value - 1)
+# write DBR (32-bit sq_head) - send doorbell at qp_dbr + 4
+q.release_mem(dbr_gpu_va + qp.qp_dbr + 4, to_be32(qp.sq_head), q.pm4.data_sel__mec_release_mem__send_32_bit_low,
+              q.pm4.int_sel__mec_release_mem__none)
+# write UAR doorbell (64-bit)
+q.release_mem(uar_gpu_va + 0x800, doorbell_val, q.pm4.data_sel__mec_release_mem__send_64_bit_data,
+              q.pm4.int_sel__mec_release_mem__none)
+q.signal(gpu.timeline_signal, gpu.next_timeline())
+q.submit(gpu)
+
+print("GPU kicked doorbell, waiting...")
+gpu.synchronize()
+
+# poll CQ from CPU (send + recv completions)
+qp.poll_cq()
+qp.poll_cq()
+
+# read back
+result = bytearray(BUF_SIZE)
+gpu.allocator._copyout(memoryview(result), buf_dst)
+gpu.synchronize()
+
+got = bytes(result[:len(test_msg)])
+print(f"result: {got}")
+assert got == test_msg, f"MISMATCH: {got} != {test_msg}"
+print("RDMA loopback send/recv test passed (GPU-kicked)")
@@ -0,0 +1,144 @@
+// MLX5 autogen header — kernel struct layouts and constants
+typedef unsigned char __u8;
+typedef unsigned short __be16;
+typedef unsigned int __be32;
+typedef unsigned long long __be64;
+
+// --- device.h structs ---
+
+struct mlx5_cmd_layout {
+    __u8    type;
+    __u8    rsvd0[3];
+    __be32  inlen;
+    __be64  in_ptr;
+    __be32  in[4];
+    __be32  out[4];
+    __be64  out_ptr;
+    __be32  outlen;
+    __u8    token;
+    __u8    sig;
+    __u8    rsvd1;
+    __u8    status_own;
+};
+
+struct mlx5_cmd_prot_block {
+    __u8    data[512];
+    __u8    rsvd0[48];
+    __be64  next;
+    __be32  block_num;
+    __u8    rsvd1;
+    __u8    token;
+    __u8    ctrl_sig;
+    __u8    sig;
+};
+
+struct mlx5_init_seg {
+    __be32  fw_rev;
+    __be32  cmdif_rev_fw_sub;
+    __be32  rsvd0[2];
+    __be32  cmdq_addr_h;
+    __be32  cmdq_addr_l_sz;
+    __be32  cmd_dbell;
+    __be32  rsvd1[120];
+    __be32  initializing;
+};
+
+// --- Command opcodes (mlx5_ifc.h) ---
+#define MLX5_CMD_OP_QUERY_HCA_CAP            0x100
+#define MLX5_CMD_OP_QUERY_ADAPTER            0x101
+#define MLX5_CMD_OP_INIT_HCA                 0x102
+#define MLX5_CMD_OP_TEARDOWN_HCA             0x103
+#define MLX5_CMD_OP_ENABLE_HCA               0x104
+#define MLX5_CMD_OP_DISABLE_HCA              0x105
+#define MLX5_CMD_OP_QUERY_PAGES              0x107
+#define MLX5_CMD_OP_MANAGE_PAGES             0x108
+#define MLX5_CMD_OP_SET_HCA_CAP              0x109
+#define MLX5_CMD_OP_QUERY_ISSI               0x10a
+#define MLX5_CMD_OP_SET_ISSI                 0x10b
+#define MLX5_CMD_OP_SET_DRIVER_VERSION       0x10d
+#define MLX5_CMD_OP_CREATE_MKEY              0x200
+#define MLX5_CMD_OP_QUERY_SPECIAL_CONTEXTS   0x203
+#define MLX5_CMD_OP_CREATE_EQ                0x301
+#define MLX5_CMD_OP_DESTROY_EQ               0x302
+#define MLX5_CMD_OP_CREATE_CQ                0x400
+#define MLX5_CMD_OP_DESTROY_CQ               0x401
+#define MLX5_CMD_OP_CREATE_QP                0x500
+#define MLX5_CMD_OP_DESTROY_QP               0x501
+#define MLX5_CMD_OP_RST2INIT_QP              0x502
+#define MLX5_CMD_OP_INIT2RTR_QP              0x503
+#define MLX5_CMD_OP_RTR2RTS_QP              0x504
+#define MLX5_CMD_OP_QUERY_NIC_VPORT_CONTEXT  0x754
+#define MLX5_CMD_OP_MODIFY_NIC_VPORT_CONTEXT 0x755
+#define MLX5_CMD_OP_SET_ROCE_ADDRESS         0x761
+#define MLX5_CMD_OP_ALLOC_PD                 0x800
+#define MLX5_CMD_OP_ALLOC_UAR                0x802
+#define MLX5_CMD_OP_ACCESS_REG               0x805
+#define MLX5_CMD_OP_ALLOC_TRANSPORT_DOMAIN   0x816
+
+// --- Command status (device.h) ---
+#define MLX5_CMD_STAT_OK                     0x0
+#define MLX5_CMD_STAT_INT_ERR                0x1
+#define MLX5_CMD_STAT_BAD_OP_ERR             0x2
+#define MLX5_CMD_STAT_BAD_PARAM_ERR          0x3
+#define MLX5_CMD_STAT_BAD_SYS_STATE_ERR      0x4
+#define MLX5_CMD_STAT_BAD_RES_ERR            0x5
+#define MLX5_CMD_STAT_RES_BUSY               0x6
+#define MLX5_CMD_STAT_LIM_ERR                0x8
+#define MLX5_CMD_STAT_BAD_RES_STATE_ERR      0x9
+#define MLX5_CMD_STAT_NO_RES_ERR             0xf
+#define MLX5_CMD_STAT_BAD_INP_LEN_ERR        0x50
+#define MLX5_CMD_STAT_BAD_OUTP_LEN_ERR       0x51
+
+// --- HCA cap types ---
+#define MLX5_CAP_GENERAL          0x0
+#define MLX5_CAP_ODP              0x2
+#define MLX5_CAP_ATOMIC           0x3
+#define MLX5_CAP_ROCE             0x4
+#define HCA_CAP_OPMOD_GET_MAX     0
+#define HCA_CAP_OPMOD_GET_CUR     1
+
+// --- Pages ---
+#define MLX5_PAGES_GIVE           1
+#define MLX5_PAGES_TAKE           2
+#define MLX5_BOOT_PAGES           1
+#define MLX5_INIT_PAGES           2
+
+// --- Registers ---
+#define MLX5_REG_HOST_ENDIANNESS  0x7004
+#define MLX5_REG_DTOR             0xC00E
+
+// --- Misc ---
+#define MLX5_PCI_CMD_XPORT        0x07
+#define MLX5_CMD_DATA_BLOCK_SIZE  512
+#define CMD_OWNER_HW              0x01
+
+// --- IFC cmd_hca_cap bit offsets ---
+#define CAP_GEN_ABS_NATIVE_PORT_NUM                       0x007
+#define CAP_GEN_HCA_CAP_2                                 0x020
+#define CAP_GEN_EVENT_ON_VHCA_STATE_ALLOCATED              0x023
+#define CAP_GEN_EVENT_ON_VHCA_STATE_ACTIVE                 0x024
+#define CAP_GEN_EVENT_ON_VHCA_STATE_IN_USE                 0x025
+#define CAP_GEN_EVENT_ON_VHCA_STATE_TEARDOWN_REQUEST       0x026
+#define CAP_GEN_LOG_MAX_QP                                 0x09B
+#define CAP_GEN_LOG_MAX_CQ                                 0x0DB
+#define CAP_GEN_RELEASE_ALL_PAGES                          0x145
+#define CAP_GEN_CACHE_LINE_128BYTE                         0x164
+#define CAP_GEN_NUM_PORTS                                  0x1B8
+#define CAP_GEN_PKEY_TABLE_SIZE                            0x190
+#define CAP_GEN_PCI_SYNC_FOR_FW_UPDATE_EVENT               0x1F1
+#define CAP_GEN_CMDIF_CHECKSUM                             0x210
+#define CAP_GEN_DCT                                        0x21A
+#define CAP_GEN_ROCE                                       0x21D
+#define CAP_GEN_ATOMIC                                     0x21E
+#define CAP_GEN_ODP                                        0x227
+#define CAP_GEN_MKEY_BY_NAME                               0x266
+#define CAP_GEN_LOG_MAX_PD                                 0x32B
+#define CAP_GEN_PCIE_RESET_USING_HOTRESET                  0x335
+#define CAP_GEN_PCI_SYNC_FOR_FW_UPDATE_WITH_DRIVER_UNLOAD  0x336
+#define CAP_GEN_VHCA_STATE                                 0x3EA
+#define CAP_GEN_ROCE_RW_SUPPORTED                          0x3A1
+#define CAP_GEN_LOG_MAX_CURRENT_UC_LIST                    0x3FB
+#define CAP_GEN_LOG_UAR_PAGE_SZ                            0x490
+#define CAP_GEN_NUM_VHCA_PORTS                             0x610
+#define CAP_GEN_SW_OWNER_ID                                0x61E
+#define CAP_GEN_NUM_TOTAL_DYNAMIC_VF_MSIX                  0x708
@@ -3,11 +3,12 @@ import os
 # TODO: there is a timing bug without this
 os.environ["AMD_AQL"] = "1"

-from tinygrad import Tensor, Device
-from tinygrad.helpers import getenv
+from tinygrad import Tensor, Device, GlobalCounters, Context
+from tinygrad.helpers import getenv, DEV
 from tinygrad.uop.ops import UOp, Ops, KernelInfo
 from tinygrad.renderer import Estimates
 from tinygrad.renderer.amd.dsl import Reg, Inst, s, v
+from tinygrad.engine.realize import run_linear

 NUM_WORKGROUPS = 96
 WAVE_SIZE = 32
@@ -36,17 +37,23 @@ def launchBenchmark(instruction, vgprIndices, dense=True, accum=False, **kwargs)
    gidx = UOp.special(NUM_WORKGROUPS, "gidx0")
    FLOPs = FLOPS_PER_MATMUL * NUM_WAVES * NUM_WORKGROUPS * INTERNAL_LOOP * INSTRUCTIONS_PER_LOOP
    sink = UOp.sink(A.base, threads, gidx, arg=KernelInfo(inst.op.name.lower(), estimates=Estimates(ops=FLOPs, mem=0)))
-    return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg="AMD"), UOp(Ops.LINEAR, src=tuple([UOp(Ops.INS, arg=x) for x in insts]))))
+    return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=Device.DEFAULT), UOp(Ops.LINEAR, src=tuple([UOp(Ops.INS, arg=x) for x in insts]))))
  dummy = Tensor.zeros(1).contiguous().realize()
  out = Tensor.custom_kernel(dummy, fxn=fxn)[0]
-  ei = out.schedule()[-1].lower()
-  elapsed = min([ei.run(wait=True) for _ in range(2)])
+  linear = out.schedule_linear()
+  ets = []
+  with Context(DEBUG=2):
+    for _ in range(2):
+      start = GlobalCounters.time_sum_s
+      run_linear(linear)
+      ets.append(GlobalCounters.time_sum_s - start)
+  elapsed = min(ets)
  FLOPs = FLOPS_PER_MATMUL * NUM_WAVES * NUM_WORKGROUPS * INTERNAL_LOOP * INSTRUCTIONS_PER_LOOP
  print(f"{inst.op_name.lower():<29} : {FLOPs/elapsed/10**12:.2f} T(FL)OPS")

 if __name__=="__main__":
  DEV = Device[Device.DEFAULT]
-  arch = DEV.renderer.arch
+  arch = DEV.renderer.target.arch

  if arch in {'gfx1100', 'gfx1103', 'gfx1151'}:
    from tinygrad.runtime.autogen.amd.rdna3.ins import *
@@ -0,0 +1,25 @@
+/* adapted from linux/drivers/gpu/drm/nouveau/include/nvfw/fw.h */
+/* SPDX-License-Identifier: MIT */
+#ifndef __NVFW_FW_H__
+#define __NVFW_FW_H__
+typedef unsigned int u32;
+
+struct nvfw_bin_hdr {
+	u32 bin_magic;
+	u32 bin_ver;
+	u32 bin_size;
+	u32 header_offset;
+	u32 data_offset;
+	u32 data_size;
+};
+
+struct nvfw_bl_desc {
+	u32 start_tag;
+	u32 dmem_load_off;
+	u32 code_off;
+	u32 code_size;
+	u32 data_off;
+	u32 data_size;
+};
+
+#endif
@@ -0,0 +1,52 @@
+/* adapted from linux/drivers/gpu/drm/nouveau/include/nvfw/hs.h */
+/* SPDX-License-Identifier: MIT */
+#ifndef __NVFW_HS_H__
+#define __NVFW_HS_H__
+typedef unsigned int u32;
+
+struct nvfw_hs_header {
+	u32 sig_dbg_offset;
+	u32 sig_dbg_size;
+	u32 sig_prod_offset;
+	u32 sig_prod_size;
+	u32 patch_loc;
+	u32 patch_sig;
+	u32 hdr_offset;
+	u32 hdr_size;
+};
+
+struct nvfw_hs_header_v2 {
+	u32 sig_prod_offset;
+	u32 sig_prod_size;
+	u32 patch_loc;
+	u32 patch_sig;
+	u32 meta_data_offset;
+	u32 meta_data_size;
+	u32 num_sig;
+	u32 header_offset;
+	u32 header_size;
+};
+
+struct nvfw_hs_load_header {
+	u32 non_sec_code_off;
+	u32 non_sec_code_size;
+	u32 data_dma_base;
+	u32 data_size;
+	u32 num_apps;
+	u32 apps[];
+};
+
+struct nvfw_hs_load_header_v2 {
+	u32 os_code_offset;
+	u32 os_code_size;
+	u32 os_data_offset;
+	u32 os_data_size;
+	u32 num_apps;
+	struct {
+		u32 offset;
+		u32 size;
+		u32 data_offset;
+		u32 data_size;
+	} app[];
+};
+#endif
@@ -10,4 +10,4 @@ def extract_ast(*args) -> None:
  return None

 if __name__ == "__main__":
-  _pmap({"get_program":extract_ast})
+  _pmap({"do_to_program":extract_ast})
@@ -4,7 +4,7 @@ import numpy as np
 from tinygrad.helpers import BEAM, Timing, CI, prod
 from tinygrad import Variable, Device, Tensor
 from tinygrad.nn import Conv2d
-from tinygrad.uop.ops import AxisType
+from tinygrad.uop.ops import AxisType, Ops
 from tinygrad.codegen.opt import Opt, OptOps
 from tinygrad.codegen.opt.postrange import Scheduler
 from tinygrad.codegen.opt.search import get_kernel_actions
@@ -84,7 +84,7 @@ class TestBeamSearch(unittest.TestCase):
    tc = Device[Device.DEFAULT].renderer.tensor_cores[0]
    size = max(tc.dims[0], tc.dims[1]) * 8
    a, b = Tensor.rand(size, size, dtype=tc.dtype_in), Tensor.rand(size, size, dtype=tc.dtype_in)
-    ast = a.matmul(b, dtype=tc.dtype_out).schedule()[-1].ast
+    ast = a.matmul(b, dtype=tc.dtype_out).schedule_linear().src[-1].src[0]
    s = Scheduler(ast, Device[Device.DEFAULT].renderer)
    s.apply_opt(Opt(OptOps.TC, 0, (-1, 0, 1)))
    up = prod([x for x, t in zip(s.full_shape, s.axis_types) if t in (AxisType.UPCAST, AxisType.UNROLL)])
@@ -94,7 +94,7 @@ class TestBeamSearch(unittest.TestCase):

  def test_max_up(self):
    a = Tensor.rand(16, 16)
-    ast = a.schedule()[-1].ast
+    ast = a.schedule_linear().src[-1].src[0]
    s = Scheduler(ast, Device[Device.DEFAULT].renderer)
    for max_up in (2, 4):
      actions = get_kernel_actions(s, include_0=False, max_up=max_up)
@@ -13,8 +13,9 @@ if __name__ == "__main__":
  devs = RemotePCIDevice.remote_list(0x1002, ((0, (0,)),), 0) or RemotePCIDevice.remote_list(0x10de, ((0, (0,)),), 0x03)
  if not devs: raise RuntimeError("no GPU found on remote")

-  pci = RemotePCIDevice("BN", devs[0])
-  print(f"connected to {os.environ['REMOTE']}, device: {devs[0]}\n")
+  sock, name = devs[0]
+  pci = RemotePCIDevice("BN", name, sock=sock)
+  print(f"connected to {os.environ['REMOTE']}, device: {name}\n")

  # ping (minimal server round-trip, no device I/O)
  from tinygrad.runtime.support.system import RemoteCmd
@@ -1,66 +0,0 @@
-# This file is automatically @generated by Cargo.
-# It is not intended for manual editing.
-version = 4
-
-[[package]]
-name = "autocfg"
-version = "1.1.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "d468802bab17cbc0cc575e9b053f41e72aa36bfa6b7f55e3529ffa43161b97fa"
-
-[[package]]
-name = "cfg-if"
-version = "1.0.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "baf1de4339761588bc0619e3cbc0120ee582ebb74b53b4efbf79117bd2da40fd"
-
-[[package]]
-name = "crunchy"
-version = "0.2.2"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "7a81dae078cea95a014a339291cec439d2f232ebe854a9d672b796c6afafa9b7"
-
-[[package]]
-name = "float-cmp"
-version = "0.9.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "98de4bbd547a563b716d8dfa9aad1cb19bfab00f4fa09a6a4ed21dbcf44ce9c4"
-dependencies = [
- "num-traits",
-]
-
-[[package]]
-name = "half"
-version = "2.3.1"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "bc52e53916c08643f1b56ec082790d1e86a32e58dc5268f897f313fbae7b4872"
-dependencies = [
- "cfg-if",
- "crunchy",
- "num-traits",
-]
-
-[[package]]
-name = "libm"
-version = "0.2.8"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "4ec2a862134d2a7d32d7983ddcdd1c4923530833c9f2ea1a44fc5fa473989058"
-
-[[package]]
-name = "num-traits"
-version = "0.2.17"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "39e3200413f237f41ab11ad6d161bc7239c84dcb631773ccd7de3dfe4b5c267c"
-dependencies = [
- "autocfg",
- "libm",
-]
-
-[[package]]
-name = "remu"
-version = "0.1.0"
-dependencies = [
- "float-cmp",
- "half",
- "num-traits",
-]
@@ -1,15 +0,0 @@
-[package]
-name = "remu"
-version = "0.1.0"
-edition = "2021"
-rust-version = "1.80.0"
-
-[lib]
-crate-type = ["cdylib"]
-
-[dependencies]
-half = { version = "2.3.1", features = ["num-traits"] }
-num-traits = "0.2.17"
-
-[dev-dependencies]
-float-cmp = "0.9.0"
@@ -1,80 +0,0 @@
-## Intro
-
-Remu is an RDNA3 emulator built to test correctness of RDNA3 code. It is used in [tinygrad's AMD CI](https://github.com/tinygrad/tinygrad).
-
-Most of the common instructions are implemented, but some formats like IMG are not supported.
-
-Remu is only for testing correctness of program output, it is not a cycle accurate simulator.
-
-## Build Locally
-
-Remu is written in Rust. Make sure you have [Cargo](https://doc.rust-lang.org/cargo/getting-started/installation.html).
-
-To build the project, run:
-
-```bash
-cargo build --release --manifest-path ./extra/remu/Cargo.toml
-```
-
-This will produce a binary in the `extra/remu/target/release` directory.
-
-## Usage with tinygrad
-
-The latest binaries are released in https://github.com/Qazalin/remu/releases. Alternatively, you can [build locally](#build-locally).
-
-Tinygrad does not yet output RDNA3 kernels directly. You can either install comgr or use `AMD_LLVM=1` (default) if you have [LLVM@19](https://github.com/tinygrad/tinygrad/blob/e2ed673c946c8f1774d816c75e52a994c2dd8a88/.github/actions/setup-tinygrad/action.yml#L208).
-
-`PYTHONPATH="." MOCKGPU=1 DEV=AMD python test/test_tiny.py TestTiny.test_plus` runs an emulated RDNA3 kernel with Remu.
-
-Add `DEBUG=6` to see Remu's logs.
-
-### DEBUG output
-
-Remu runs each thread one at a time in a nested for loop, see lib.rs. The DEBUG output prints information about the current thread.
-
-The DEBUG output has 3 sections:
-
-```
-<------------ 1 ----------> <--- 2 ---> <--------------------------------------- 3 ------------------------------------------>
-[0   0   0  ] [0   0   0  ] 0  F4080100 SMEM { op: 2, sdata: 4, sbase: 0, offset: 0, soffset: 124, glc: false, dlc: false }
-```
-
-#### Section 1: Grid info
-
-`[gid.x, gid.y, gid.z], [lid.x, lid.y, lid.z]` of the current thread.
-
-#### Section 2: Wave info
-
-`<lane> <instruction hex>`
-
-RDNA3 divides threads into chunks of 32. Each thread is assigned to a "lane" from 0-31.
-
-In Remu, even though all threads run one at a time, each 32 thread chunk (a wave) shares state like SGPR, VGPR, LDS, EXEC mask, etc.
-Remu can simulate up to one wave sync instruction.
-For more details, see work_group.rs.
-
-Section 2 can have a green or gray color.
-
-Green = The thread is actively executing the instruction.
-
-Gray = The thread has been "turned off" by the EXEC mask, it skips execution of some instructions. (refer to "EXECute Mask" on [page 23](https://www.amd.com/content/dam/amd/en/documents/radeon-tech-docs/instruction-set-architectures/rdna3-shader-instruction-set-architecture-feb-2023_0.pdf#page=23) of ISA docs for more details.)
-
-To see the colors in action, try running `DEBUG=6 PYTHONPATH="." MOCKGPU=1 DEV=AMD python test/test_ops.py TestOps.test_arange_big`. See how only lane 0 writes to global memory:
-```
-[255 0   0  ] [0   0   0  ] 0  DC6A0000 FLAT { op: 26, offset: 0, dlc: false, glc: false, slc: false, seg: 2, addr: 8, data: 0, saddr: 0, sve: false, vdst: 0 }
-[255 0   0  ] [1   0   0  ] 1  DC6A0000
-[255 0   0  ] [2   0   0  ] 2  DC6A0000
-[255 0   0  ] [3   0   0  ] 3  DC6A0000
-[255 0   0  ] [3   0   0  ] 4  DC6A0000
-```
-
-#### Section 3: Decoded Instruction
-
-This prints the instruction type and all the parsed bitfields.
-
-Remu output vs llvm-objdump:
-
-```
-s_load_b64 s[0:1], s[0:1], 0x10                            // 00000000160C: F4040000 F8000010
-SMEM { op: 1, sdata: 0, sbase: 0, offset: 16, soffset: 124, glc: false, dlc: false }
-```
@@ -1 +0,0 @@
-max_width = 150
@@ -1,162 +0,0 @@
-use half::f16;
-use num_traits::{float::FloatCore, PrimInt, Unsigned, clamp};
-
-pub fn bits<T>(word: T, hi: usize, lo: usize) -> T where T: PrimInt + Unsigned {
-    assert!(hi >= lo);
-    let width = hi - lo + 1;
-    (word >> lo) & ((T::one() << width) - T::one())
-}
-
-pub fn nth(val: u32, pos: usize) -> u32 {
-    (val >> (31 - pos as u32)) & 1
-}
-pub fn f16_lo(val: u32) -> f16 {
-    f16::from_bits((val & 0xffff) as u16)
-}
-pub fn f16_hi(val: u32) -> f16 {
-    f16::from_bits(((val >> 16) & 0xffff) as u16)
-}
-
-pub fn sign_ext(num: u64, bits: usize) -> i64 {
-    let mut value = num;
-    let is_negative = (value >> (bits - 1)) & 1 != 0;
-    if is_negative {
-        value |= !0 << bits;
-    }
-    value as i64
-}
-
-pub trait IEEEClass<T> {
-    fn exponent(&self) -> T;
-}
-impl IEEEClass<u32> for f32 {
-    fn exponent(&self) -> u32 {
-        (self.to_bits() & 0b01111111100000000000000000000000) >> 23
-    }
-}
-impl IEEEClass<u16> for f16 {
-    fn exponent(&self) -> u16 {
-        (self.to_bits() & 0b0111110000000000) >> 10
-    }
-}
-impl IEEEClass<u64> for f64 {
-    fn exponent(&self) -> u64 {
-        (self.to_bits() & 0b0111111111110000000000000000000000000000000000000000000000000000) >> 52
-    }
-}
-
-pub trait VOPModifier<T> {
-    fn negate(&self, pos: usize, modifier: usize) -> T;
-    fn absolute(&self, pos: usize, modifier: usize) -> T;
-    fn clmp(&self, cm: bool) -> T;
-}
-impl<T> VOPModifier<T> for T
-where
-    T: FloatCore,
-{
-    fn negate(&self, pos: usize, modifier: usize) -> T {
-        match (modifier >> pos) & 1 {
-            1 => -*self,
-            _ => *self,
-        }
-    }
-    fn absolute(&self, pos: usize, modifier: usize) -> T {
-        match (modifier >> pos) & 1 {
-            1 => self.abs(),
-            _ => *self,
-        }
-    }
-    fn clmp(&self, cm:bool) -> T {
-        if !cm { return *self }
-        let r = clamp(*self, T::zero(), T::one());
-        if r == T::zero() { T::zero() } else { r }
-    }
-}
-
-pub fn extract_mantissa(x: f64) -> f64 {
-    if x.is_infinite() || x.is_nan() {
-        return x;
-    }
-    let bits = x.to_bits();
-    let mantissa_mask: u64 = 0x000FFFFFFFFFFFFF;
-    let bias: u64 = 1023;
-    let normalized_mantissa_bits = (bits & mantissa_mask) | ((bias - 1) << 52);
-    return f64::from_bits(normalized_mantissa_bits);
-}
-pub fn ldexp(x: f64, exp: i32) -> f64 {
-    x * 2f64.powi(exp)
-}
-
-#[cfg(test)]
-mod tests {
-    use super::*;
-    #[test]
-    fn test_extract_mantissa() {
-        assert_eq!(extract_mantissa(2.0f64), 0.5);
-    }
-
-    #[test]
-    fn test_normal_exponent() {
-        assert_eq!(2.5f32.exponent(), 128);
-        assert_eq!(1.17549435e-38f32.exponent(), 1);
-        assert_eq!(f32::INFINITY.exponent(), 255);
-        assert_eq!(f32::NEG_INFINITY.exponent(), 255);
-    }
-
-    #[test]
-    fn test_denormal_exponent() {
-        assert_eq!(1.0e-40f32.exponent(), 0);
-        assert_eq!(1.0e-42f32.exponent(), 0);
-        assert_eq!(1.0e-44f32.exponent(), 0);
-        assert_eq!((1.17549435e-38f32 / 2.0).exponent(), 0);
-    }
-
-    #[test]
-    fn test_normal_exponent_f16() {
-        assert_eq!(f16::from_f32(3.14f32).exponent(), 16);
-        assert_eq!(f16::NEG_INFINITY.exponent(), 31);
-        assert_eq!(f16::INFINITY.exponent(), 31);
-    }
-
-    #[test]
-    fn test_neg() {
-        assert_eq!(0.3_f32.negate(0, 0b001), -0.3_f32);
-        assert_eq!(0.3_f32.negate(1, 0b010), -0.3_f32);
-        assert_eq!(0.3_f32.negate(2, 0b100), -0.3_f32);
-        assert_eq!(0.3_f32.negate(0, 0b110), 0.3_f32);
-        assert_eq!(0.3_f32.negate(1, 0b010), -0.3_f32);
-        assert_eq!(0.0_f32.negate(0, 0b001).to_bits(), (-0.0f32).to_bits());
-        assert_eq!((-0.0_f32).negate(0, 0b001).to_bits(), 0);
-    }
-
-    #[test]
-    fn test_sign_ext() {
-        assert_eq!(sign_ext(0b000000000000000101000, 21), 40);
-        assert_eq!(sign_ext(0b111111111111111011000, 21), -40);
-        assert_eq!(sign_ext(0b000000000000000000000, 21), 0);
-        assert_eq!(sign_ext(0b111111111111111111111, 21), -1);
-        assert_eq!(sign_ext(0b111000000000000000000, 21), -262144);
-        assert_eq!(sign_ext(0b000111111111111111111, 21), 262143);
-        assert_eq!(sign_ext(7608, 13), -584);
-    }
-}
-
-use std::sync::LazyLock;
-pub static DEBUG: LazyLock<bool> = LazyLock::new(|| std::env::var("DEBUG").map(|v| v.parse::<usize>().unwrap_or(0) >= 6).unwrap_or(false));
-
-pub fn colored(st:&str, color:&str) -> String {
-    let ansi_code = match color {
-        "green" => format!("\x1b[{};2;39;176;139m", 38),
-        "gray" => format!("\x1b[{};2;169;169;169m", 38),
-        _ => format!("\x1b[{};2;255;255;255m", 38),
-    };
-    format!("{}{}{}", ansi_code, st, "\x1b[0m")
-}
-
-#[macro_export]
-macro_rules! todo_instr {
-    ($x:expr) => {{
-        println!("{:08X}", $x);
-        Err(1)
-    }};
-}
@@ -1,77 +0,0 @@
-use crate::state::StateSnapshot;
-use crate::work_group::{WaveContext, WorkGroup};
-use std::os::raw::c_char;
-use std::slice;
-mod helpers;
-mod rdna3;
-mod state;
-mod thread;
-mod work_group;
-
-#[no_mangle]
-pub extern "C" fn run_asm(lib: *const c_char, lib_sz: u32, gx: u32, gy: u32, gz: u32, lx: u32, ly: u32, lz: u32, args_ptr: *const u64) -> i32 {
-    if lib.is_null() || (lib_sz % 4) != 0 {
-        panic!("Pointer is null or length is not properly aligned to 4 bytes");
-    }
-    let kernel = unsafe { slice::from_raw_parts(lib as *const u32, (lib_sz / 4) as usize).to_vec() };
-    let dispatch_dim = match (gy != 1, gz != 1) {
-        (true, true) => 3,
-        (true, false) => 2,
-        _ => 1,
-    };
-    for gx in 0..gx {
-        for gy in 0..gy {
-            for gz in 0..gz {
-                let mut wg = WorkGroup::new(dispatch_dim, [gx, gy, gz], [lx, ly, lz], &kernel, args_ptr);
-                if let Err(err) = wg.exec_waves() {
-                    return err;
-                }
-            }
-        }
-    }
-    0
-}
-
-// FFI functions for single-stepping comparison tests
-
-#[no_mangle]
-pub extern "C" fn wave_create(lib: *const c_char, lib_sz: u32, n_lanes: u32) -> *mut WaveContext {
-    if lib.is_null() || (lib_sz % 4) != 0 { return std::ptr::null_mut(); }
-    let kernel = unsafe { slice::from_raw_parts(lib as *const u32, (lib_sz / 4) as usize).to_vec() };
-    Box::into_raw(Box::new(WaveContext::new(kernel, n_lanes as usize)))
-}
-
-#[no_mangle]
-pub extern "C" fn wave_step(ctx: *mut WaveContext) -> i32 {
-    if ctx.is_null() { return -99; }
-    unsafe { (*ctx).step() }
-}
-
-#[no_mangle]
-pub extern "C" fn wave_get_snapshot(ctx: *const WaveContext, out: *mut StateSnapshot) {
-    if ctx.is_null() || out.is_null() { return; }
-    unsafe { *out = (*ctx).get_snapshot(); }
-}
-
-#[no_mangle]
-pub extern "C" fn wave_set_sgpr(ctx: *mut WaveContext, idx: u32, val: u32) {
-    if ctx.is_null() || idx >= 128 { return; }
-    unsafe { (*ctx).scalar_reg[idx as usize] = val; }
-}
-
-#[no_mangle]
-pub extern "C" fn wave_set_vgpr(ctx: *mut WaveContext, lane: u32, idx: u32, val: u32) {
-    if ctx.is_null() || lane >= 32 || idx >= 256 { return; }
-    unsafe { (*ctx).vec_reg.get_lane_mut(lane as usize)[idx as usize] = val; }
-}
-
-#[no_mangle]
-pub extern "C" fn wave_init_lds(ctx: *mut WaveContext, size: u32) {
-    if ctx.is_null() { return; }
-    unsafe { (*ctx).lds.data.resize(size as usize, 0); }
-}
-
-#[no_mangle]
-pub extern "C" fn wave_free(ctx: *mut WaveContext) {
-    if !ctx.is_null() { unsafe { drop(Box::from_raw(ctx)); } }
-}
@@ -1,223 +0,0 @@
-use crate::helpers::{bits, sign_ext};
-
-#[derive(Debug, PartialEq)]
-pub enum Instruction {
-    SOP2 { op: u8, ssrc0: u8, ssrc1: u8, sdst: u8 },
-    SOP1 { op: u8, ssrc0: u8, sdst: u8 },
-    SOPK { op: u8, simm16: i16, sdst: u8 },
-    SOPP { op: u8, simm16: i16 },
-    SOPC { op: u8, ssrc0: u8, ssrc1: u8 },
-
-    SMEM { op: u8, sdata: u8, sbase: u8, offset: i32, soffset: u8, glc: bool, dlc: bool },
-
-    VOP1 { op: u8, vdst: u8, src: u16 },
-    VOP2 { op: u8, vdst: u8, vsrc: u8, src: u16 },
-    VOPC { op: u8, vsrc: u8, src: u16 },
-    VOP3 { op: u32, opsel: u8, cm: bool, abs: u8, vdst: u8, neg: u8, omod: u8, src2: u16, src1: u16, src0: u16 },
-    VOP3SD { op: u32, cm: bool, sdst: u8, vdst: u8, neg: u8, omod: u8, src2: u16, src1: u16, src0: u16 },
-    VOP3P { op: u8, vdst: u8, neg_hi: u8, opsel: u8, opsel_hi: u8, opsel_hi2: bool, cm: bool, src2: u16, src1: u16, src0: u16, neg: u8 },
-    VOPD { opx: u8, opy: u8, vdstx: u8, vdsty: u8, vsrcx1: u8, vsrcy1: u8, srcx0: u16, srcy0: u16 },
-
-    DS { op: u8, gds: bool, offset1: u8, offset0: u8, vdst: u8, data1: u8, data0: u8, addr: u8 },
-
-    FLAT { op: u8, offset: u16, dlc: bool, glc: bool, slc: bool, seg: u8, addr: u8, data: u8, saddr: u8, sve: bool, vdst: u8 }
-}
-
-const VOP3SD_OPS: [u32; 7] = [764, 765, 766, 767, 768, 769, 770];
-
-pub fn decode(word:u32, word1:Option<&u32>) -> Instruction {
-    match bits(word, 31, 30) {
-        0b11 => {
-            let word = (*word1.unwrap() as u64) << 32 | (word as u64);
-            match bits(word, 29, 26) {
-                0b1101 => {
-                    let sbase = (bits(word, 5, 0) as u8) << 1;
-                    let sdata = bits(word, 12, 6) as u8;
-                    let dlc = bits(word, 13, 13) != 0;
-                    let glc = bits(word, 14, 14) != 0;
-                    let op = bits(word, 25, 18) as u8;
-                    let offset = sign_ext(bits(word, 52, 32), 21) as i32;
-                    let soffset = bits(word, 63, 57) as u8;
-                    Instruction::SMEM { sbase, sdata, dlc, glc, op, offset, soffset }
-                }
-                0b0101 => {
-                    let op = bits(word, 25, 16) as u32;
-                    let vdst = bits(word, 7, 0) as u8;
-                    let cm = bits(word, 15, 15) != 0;
-                    let src0 = bits(word, 40, 32) as u16;
-                    let src1 = bits(word, 49, 41) as u16;
-                    let src2 = bits(word, 58, 50) as u16;
-                    let omod = bits(word, 60, 59) as u8;
-                    let neg = bits(word, 63, 61) as u8;
-                     if VOP3SD_OPS.contains(&op) {
-                         let sdst = bits(word, 14, 8) as u8;
-                         Instruction::VOP3SD { op, vdst, sdst, cm, src0, src1, src2, omod, neg }
-                     } else {
-                        let abs = bits(word, 10, 8) as u8;
-                        let opsel = bits(word, 14, 11) as u8;
-                        Instruction::VOP3 { opsel, cm, abs, vdst, neg, omod, src2, src1, src0, op }
-                     }
-                }
-                0b0011 => {
-                    let op = bits(word, 22, 16) as u8;
-                    let vdst = bits(word, 7, 0) as u8;
-                    let neg_hi = bits(word, 10, 8) as u8;
-                    let opsel = bits(word, 13, 11) as u8;
-                    let opsel_hi2 = bits(word, 14, 14) != 0;
-                    let cm = bits(word, 15, 15) != 0;
-                    let src0 = bits(word, 40, 32) as u16;
-                    let src1 = bits(word, 49, 41) as u16;
-                    let src2 = bits(word, 58, 50) as u16;
-                    let opsel_hi = bits(word, 60, 59) as u8;
-                    let neg = bits(word, 63, 61) as u8;
-                    Instruction::VOP3P { op, vdst, neg_hi, opsel, opsel_hi, opsel_hi2, cm, src0, src1, src2, neg }
-                }
-                0b0110 => {
-                    let offset0 = bits(word, 7, 0) as u8;
-                    let offset1 = bits(word, 15, 8) as u8;
-                    let gds = bits(word, 17, 17) != 0;
-                    let op = bits(word, 25, 18) as u8;
-                    let addr = bits(word, 39, 32) as u8;
-                    let data0 = bits(word, 47, 40) as u8;
-                    let data1 = bits(word, 55, 48) as u8;
-                    let vdst = bits(word, 63, 56) as u8;
-                    Instruction::DS { op, gds, offset1, offset0, vdst, data1, data0, addr }
-                }
-                0b0111 => {
-                    let offset = bits(word, 12, 0) as u16;
-                    let dlc = bits(word, 13, 13) != 0;
-                    let glc = bits(word, 14, 14) != 0;
-                    let slc = bits(word, 15, 15) != 0;
-                    let seg = bits(word, 17, 16) as u8;
-                    let op = bits(word, 24, 18) as u8;
-                    let addr = bits(word, 39, 32) as u8;
-                    let data = bits(word, 47, 40) as u8;
-                    let saddr = bits(word, 54, 48) as u8;
-                    let sve = bits(word, 55, 55) != 0;
-                    let vdst = bits(word, 63, 56) as u8;
-                    Instruction::FLAT { offset, dlc, glc, slc, seg, op, addr, data, saddr, sve, vdst }
-                },
-                0b0010 => {
-                    let srcx0 = bits(word, 8, 0) as u16;
-                    let vsrcx1 = bits(word, 16, 9) as u8;
-                    let opy = bits(word, 21, 17) as u8;
-                    let opx = bits(word, 25, 22) as u8;
-                    let srcy0 = bits(word, 40, 32) as u16;
-                    let vsrcy1 = bits(word, 48, 41) as u8;
-                    let vdsty = bits(word, 55, 49) as u8;
-                    let vdstx = bits(word, 63, 56) as u8;
-                    Instruction::VOPD { opx, opy, vdstx, vdsty, vsrcx1, vsrcy1, srcx0, srcy0 }
-                }
-                _ => todo!(),
-            }
-        }
-        0b10 => {
-            let ssrc0 = bits(word, 7, 0) as u8;
-            let ssrc1 = bits(word, 15, 8) as u8;
-            let simm16 = word as i16;
-            let sdst = bits(word, 22, 16) as u8;
-            match bits(word, 29, 23) {
-                0b1111101 => Instruction::SOP1 { ssrc0, sdst, op: bits(word, 15, 8) as u8 },
-                0b1111110 => Instruction::SOPC { ssrc0, ssrc1, op: bits(word, 22, 16) as u8 },
-                0b1111111 => Instruction::SOPP { simm16, op: bits(word, 22, 16) as u8 },
-                _ => {
-                    match bits(word, 29, 28) {
-                        0b11 => Instruction::SOPK { simm16, sdst, op: bits(word, 27, 23) as u8 },
-                        _ => Instruction::SOP2 { ssrc0, ssrc1, sdst, op: bits(word, 29, 23) as u8 }
-                    }
-                }
-            }
-        }
-        _ => {
-            let vdst = bits(word, 24, 17) as u8;
-            let src = bits(word, 8, 0) as u16;
-            let vsrc = bits(word, 16, 9) as u8;
-            match bits(word, 30, 25) {
-                0b111110 => Instruction::VOPC { vsrc, src, op: bits(word, 24, 17) as u8 },
-                0b111111 => Instruction::VOP1 { vdst, src, op: vsrc },
-                _ => Instruction::VOP2 { vdst, vsrc, src, op: bits(word, 30, 25) as u8 },
-            }
-        },
-    }
-}
-
-#[cfg(test)]
-mod test_rdna3 {
-    use super::*;
-
-    use std::process::{Stdio, Command};
-    use std::io::{Result, Write};
-
-    const LLVM_ARGS: &[&str; 3] = &["--arch=amdgcn", "--mcpu=gfx1100", "--triple=amdgcn-amd-amdhsa"];
-    const OFFSET_PRG: usize = 16;
-    const NULL: u8 = 124;
-
-    fn llvm_assemble(asm: &str) -> Result<Vec<u8>> {
-        let mut proc = Command::new("llvm-mc").args(LLVM_ARGS).args(["-filetype=obj", "-o", "-"]).stdin(Stdio::piped()).stdout(Stdio::piped()).spawn()?;
-        proc.stdin.as_mut().unwrap().write_all(asm.as_bytes())?;
-        let out = proc.wait_with_output()?;
-        match out.status.success() {
-            true => Ok(out.stdout),
-            false => Err(std::io::Error::new(std::io::ErrorKind::Other, "llvm-mc err")),
-        }
-    }
-
-    fn llvm_disassemble(code: &Vec<u8>) -> Result<String> {
-        let mut proc = Command::new("llvm-objdump").args(LLVM_ARGS).args(["--disassemble", "-"]).stdin(Stdio::piped()).stdout(Stdio::piped()).spawn()?;
-        proc.stdin.as_mut().unwrap().write_all(code)?;
-        let out = proc.wait_with_output()?;
-        match out.status.success() {
-            true => Ok(String::from_utf8(out.stdout).unwrap()),
-            false => Err(std::io::Error::new(std::io::ErrorKind::Other, "llvm-objdump err")),
-        }
-    }
-
-    fn test_decode(asm: &str) -> Instruction {
-        let lib = llvm_assemble(asm).unwrap();
-        println!("{}", llvm_disassemble(&lib).unwrap());
-        let stream: Vec<u32> = lib.chunks_exact(4).map(|chunk| u32::from_le_bytes(chunk.try_into().unwrap())).skip(OFFSET_PRG).collect();
-        decode(stream[0], stream.get(1))
-    }
-
-    #[test]
-    fn test_decode_smem() {
-        assert_eq!(test_decode("s_load_b128 s[4:7], s[0:1], null"), Instruction::SMEM { op: 2, sdata: 4, sbase: 0, offset: 0, soffset: NULL, glc: false, dlc: false });
-        assert_eq!(test_decode("s_load_b32 s10, s[0:1], 0xc"), Instruction::SMEM { op: 0, sdata: 10, sbase: 0, offset: 0xc, soffset: NULL, glc: false, dlc: false });
-        assert_eq!(test_decode("s_load_b32 s0, s[4:5], s6"), Instruction::SMEM { op: 0, sdata: 0, sbase: 4, offset: 0, soffset: 6, glc: false, dlc: false });
-        assert_eq!(test_decode("s_load_b32 s0, s[4:5], glc dlc"), Instruction::SMEM { op: 0, sdata: 0, sbase: 4, offset: 0, soffset: NULL, glc: true, dlc: true });
-        assert_eq!(test_decode("s_load_b32 s0, s[4:5], glc"), Instruction::SMEM { op: 0, sdata: 0, sbase: 4, offset: 0, soffset: NULL, glc: true, dlc: false });
-        assert_eq!(test_decode("s_load_b32 s0, s[4:5], -20"), Instruction::SMEM { op: 0, sdata: 0, sbase: 4, offset: -20, soffset: NULL, glc: false, dlc: false });
-        assert_eq!(test_decode("s_load_b32 s0, s[4:5], -1048576"), Instruction::SMEM { op: 0, sdata: 0, sbase: 4, offset: -1048576, soffset: NULL, glc: false, dlc: false });
-    }
-
-    #[test]
-    fn test_decode_salu() {
-        assert_eq!(test_decode("s_add_u32 s1 s2 s3"), Instruction::SOP2 { op: 0, ssrc0: 2, ssrc1: 3, sdst: 1 });
-        assert_eq!(test_decode("s_add_u32 vcc_hi exec_lo vcc_lo"), Instruction::SOP2 { op: 0, ssrc0: 126, ssrc1: 106, sdst: 107 });
-        assert_eq!(test_decode("s_mov_b32 s1 -0.5"), Instruction::SOP1 { op: 0, ssrc0: 241, sdst: 1 });
-        assert_eq!(test_decode("s_cmpk_eq_i32 s0 -30"), Instruction::SOPK { op: 3, sdst: 0, simm16: -30 });
-        assert_eq!(test_decode("s_cmpk_eq_u32 s0 65535"), Instruction::SOPK { op: 9, sdst: 0, simm16: -1 });
-        assert_eq!(test_decode("s_cmp_ge_i32 s1 s2"), Instruction::SOPC { op: 3, ssrc0: 1, ssrc1: 2 });
-    }
-
-    #[test]
-    fn test_decode_valu_e32() {
-        assert_eq!(test_decode("v_mov_b32 v0, v0"), Instruction::VOP1 { op: 1, vdst: 0, src: 256 });
-        assert_eq!(test_decode("v_mov_b32 v0, s0"), Instruction::VOP1 { op: 1, vdst: 0, src: 0 });
-        assert_eq!(test_decode("v_cmp_t_f32 v1, v0"), Instruction::VOPC { op: 31, vsrc: 0, src: 257 });
-    }
-
-    #[test]
-    fn test_decode_valu_e64() {
-        assert_eq!(test_decode("v_log_f32_e64 v2, |v0|"), Instruction::VOP3 { op: 423, vdst: 2, src0: 256, src1: 0, src2: 0, abs: 0b001, neg: 0, opsel: 0, omod: 0, cm: false });
-        assert_eq!(test_decode("v_div_scale_f32 v2, s1, v0, v1, v2"), Instruction::VOP3SD { op: 764, cm: false, vdst: 2, sdst: 1, src0: 256, src1: 257, src2: 258, omod: 0, neg: 0 });
-        assert_eq!(test_decode("v_pk_add_i16 v1, v0, v2"), Instruction::VOP3P { op: 2, vdst: 1, neg_hi: 0, opsel: 0, opsel_hi: 3, opsel_hi2: true, cm: false, src2: 0, src1: 258, src0: 256, neg: 0 });
-    }
-
-    #[test]
-    fn test_decode_ds() {
-        assert_eq!(test_decode("ds_add_u32 v2, v4 offset:16"), Instruction::DS { op: 0, gds: false, offset1: 0, offset0: 0x10, vdst: 0, data1: 0, data0: 4, addr: 2 });
-        assert_eq!(test_decode("ds_store_b32 v0, v1, offset: 0x04 gds"), Instruction::DS { op: 13, gds: true, offset1: 0, offset0: 0x04, vdst: 0, data1: 0, data0: 1, addr: 0 });
-        assert_eq!(test_decode("ds_load_u8 v1, v0 offset:16"), Instruction::DS { op: 58, gds: false, offset1: 0, offset0: 16, vdst: 1, data1: 0, data0: 0, addr: 0 });
-    }
-}
@@ -1,272 +0,0 @@
-use std::ops::{Index, IndexMut};
-
-pub trait Register {
-    fn read64(&self, idx: usize) -> u64;
-    fn write64(&mut self, idx: usize, addr: u64);
-}
-impl<T> Register for T where T: Index<usize, Output = u32> + IndexMut<usize> {
-    fn read64(&self, idx: usize) -> u64 {
-        let lsb = self[idx] as u64;
-        let msb = self[idx + 1] as u64;
-        (msb << 32) | lsb
-    }
-
-    fn write64(&mut self, idx: usize, value: u64) {
-        self[idx] = (value & 0xffffffff) as u32;
-        self[idx + 1] = ((value & (0xffffffff << 32)) >> 32) as u32;
-    }
-}
-
-#[derive(Debug, Clone)]
-pub struct VGPR {
-    values: [[u32; 256]; 32],
-    pub default_lane: Option<usize>,
-}
-impl Index<usize> for VGPR {
-    type Output = u32;
-    fn index(&self, index: usize) -> &Self::Output {
-        &self.values[self.default_lane.unwrap()][index]
-    }
-}
-impl IndexMut<usize> for VGPR {
-    fn index_mut(&mut self, index: usize) -> &mut Self::Output {
-        &mut self.values[self.default_lane.unwrap()][index]
-    }
-}
-impl VGPR {
-    pub fn new() -> Self {
-        VGPR {
-            values: [[0; 256]; 32],
-            default_lane: None,
-        }
-    }
-    pub fn get_lane(&self, lane: usize) -> [u32; 256] {
-        *self.values.get(lane).unwrap()
-    }
-    pub fn get_lane_mut(&mut self, lane: usize) -> &mut [u32; 256] {
-        self.values.get_mut(lane).unwrap()
-    }
-}
-
-pub trait Value {
-    fn mut_hi16(&mut self, val: u16);
-    fn mut_lo16(&mut self, val: u16);
-}
-impl Value for u32 {
-    fn mut_hi16(&mut self, val: u16) {
-        *self = ((val as u32) << 16) | (*self as u16 as u32);
-    }
-    fn mut_lo16(&mut self, val: u16) {
-        *self = ((((*self & (0xffff << 16)) >> 16) as u32) << 16) | val as u32;
-    }
-}
-
-#[derive(Debug, Clone, Copy)]
-pub struct WaveValue {
-    pub value: u32,
-    pub warp_size: usize,
-    pub default_lane: Option<usize>,
-    pub mutations: Option<[bool; 32]>,
-}
-impl WaveValue {
-    pub fn new(value: u32, warp_size: usize) -> Self {
-        Self {
-            value,
-            warp_size,
-            default_lane: None,
-            mutations: None,
-        }
-    }
-    pub fn read(&self) -> bool {
-        (self.value >> self.default_lane.unwrap()) & 1 == 1
-    }
-    pub fn set_lane(&mut self, value: bool) {
-        if self.mutations.is_none() {
-            self.mutations = Some([false; 32])
-        }
-        self.mutations.as_mut().unwrap()[self.default_lane.unwrap()] = value;
-    }
-    pub fn apply_muts(&mut self) {
-        self.value = 0;
-        for lane in 0..self.warp_size {
-            if self.mutations.unwrap()[lane] {
-                self.value |= 1 << lane;
-            }
-        }
-    }
-}
-
-/// C-compatible state snapshot for FFI - used for comparing emulator states
-#[repr(C)]
-#[derive(Clone, Debug)]
-pub struct StateSnapshot {
-    pub pc: u32,
-    pub scc: u32,
-    pub vcc: u32,
-    pub exec_mask: u32,
-    pub sgpr: [u32; 128],
-    pub vgpr: [[u32; 256]; 32],
-}
-
-impl StateSnapshot {
-    pub fn new() -> Self {
-        Self { pc: 0, scc: 0, vcc: 0, exec_mask: 0, sgpr: [0; 128], vgpr: [[0; 256]; 32] }
-    }
-}
-
-#[derive(Clone, Debug)]
-pub struct VecDataStore {
-    pub data: Vec<u8>,
-}
-
-impl VecDataStore {
-    pub fn new() -> Self {
-        Self { data: Vec::new() }
-    }
-    pub fn write(&mut self, addr: usize, val: u32) {
-        if addr + 4 >= self.data.len() {
-            self.data.resize(self.data.len() + addr + 5, 0);
-        }
-        self.data[addr..addr + 4].iter_mut().enumerate().for_each(|(i, x)| {
-            *x = val.to_le_bytes()[i];
-        });
-    }
-    pub fn write64(&mut self, addr: usize, val: u64) {
-        self.write(addr, (val & 0xffffffff) as u32);
-        self.write(addr + 4, ((val & (0xffffffff << 32)) >> 32) as u32);
-    }
-    pub fn read(&self, addr: usize) -> u32 {
-        let mut bytes: [u8; 4] = [0; 4];
-        bytes.copy_from_slice(&self.data[addr + 0..addr + 4]);
-        u32::from_le_bytes(bytes)
-    }
-    pub fn read64(&mut self, addr: usize) -> u64 {
-        let lsb = self.read(addr);
-        let msb = self.read(addr + 4);
-        ((msb as u64) << 32) | lsb as u64
-    }
-}
-
-#[cfg(test)]
-mod test_state {
-    use super::*;
-
-    #[test]
-    fn test_wave_value() {
-        let mut val = WaveValue::new(0b11000000000000011111111111101110, 32);
-        val.default_lane = Some(0);
-        assert!(!val.read());
-        val.default_lane = Some(31);
-        assert!(val.read());
-    }
-
-    #[test]
-    fn test_wave_value_small() {
-        let mut val = WaveValue::new(0, 1);
-        val.default_lane = Some(0);
-        assert!(!val.read());
-        assert_eq!(val.value, 0);
-        val.set_lane(true);
-        val.apply_muts();
-        assert!(val.read());
-        assert_eq!(val.value, 1);
-    }
-
-    #[test]
-    fn test_wave_value_small_alt() {
-        let mut val = WaveValue::new(0, 2);
-        val.default_lane = Some(0);
-        assert!(!val.read());
-        assert_eq!(val.value, 0);
-        val.set_lane(true);
-        val.apply_muts();
-        assert!(val.read());
-        assert_eq!(val.value, 1);
-    }
-
-    #[test]
-    fn test_wave_value_exec() {
-        let warp_size = 32;
-        let val = WaveValue::new(u32::MAX, warp_size);
-        assert_eq!(val.value, u32::MAX);
-        let warp_size = 3;
-        let val = WaveValue::new((1 << warp_size) - 1, warp_size);
-        assert_eq!(val.value, 7)
-    }
-
-    #[test]
-    fn test_wave_value_toggle_one() {
-        let warp_size = 2;
-        let mut val = WaveValue::new(0b11, warp_size);
-        // 0
-        val.default_lane = Some(0);
-        val.set_lane(false);
-        // 1
-        val.default_lane = Some(1);
-        val.set_lane(true);
-        val.apply_muts();
-        assert_eq!(val.value, 2);
-    }
-
-    #[test]
-    fn test_wave_value_mutate_small() {
-        let mut val = WaveValue::new(0, 2);
-        val.default_lane = Some(0);
-        assert!(!val.read());
-        assert_eq!(val.value, 0);
-        val.set_lane(true);
-        val.apply_muts();
-        assert!(val.read());
-        assert_eq!(val.value, 1);
-    }
-
-    #[test]
-    fn test_wave_value_mutations() {
-        let mut val = WaveValue::new(0b10001, 32);
-        val.default_lane = Some(0);
-        val.set_lane(false);
-        assert!(val.mutations.unwrap().iter().all(|x| !x));
-        val.default_lane = Some(1);
-        val.set_lane(true);
-        assert_eq!(val.value, 0b10001);
-        assert_eq!(
-            val.mutations,
-            Some([
-                false, true, false, false, false, false, false, false, false, false, false, false, false, false, false, false, false, false, false,
-                false, false, false, false, false, false, false, false, false, false, false, false, false,
-            ])
-        );
-
-        val.apply_muts();
-        assert_eq!(val.value, 0b10);
-    }
-
-    #[test]
-    fn test_write16() {
-        let mut vgpr = VGPR::new();
-        vgpr.default_lane = Some(0);
-        vgpr[0] = 0b11100000000000001111111111111111;
-        vgpr[0].mut_lo16(0b1011101111111110);
-        assert_eq!(vgpr[0], 0b11100000000000001011101111111110);
-    }
-
-    #[test]
-    fn test_write16hi() {
-        let mut vgpr = VGPR::new();
-        vgpr.default_lane = Some(0);
-        vgpr[0] = 0b11100000000000001111111111111111;
-        vgpr[0].mut_hi16(0b1011101111111110);
-        assert_eq!(vgpr[0], 0b10111011111111101111111111111111);
-    }
-
-    #[test]
-    fn test_vgpr() {
-        let mut vgpr = VGPR::new();
-        vgpr.default_lane = Some(0);
-        vgpr[0] = 42;
-        vgpr.default_lane = Some(10);
-        vgpr[0] = 10;
-        assert_eq!(vgpr.get_lane(0)[0], 42);
-        assert_eq!(vgpr.get_lane(10)[0], 10);
-    }
-}
@@ -1,323 +0,0 @@
-use crate::helpers::{colored, DEBUG};
-use crate::state::{Register, StateSnapshot, VecDataStore, WaveValue, VGPR};
-use crate::thread::{Thread, END_PRG, SGPR_COUNT};
-use std::collections::HashMap;
-
-pub const WAVE_SIZE: usize = 32;
-
-pub struct WorkGroup<'a> {
-    dispatch_dim: u32,
-    id: [u32; 3],
-    lds: VecDataStore,
-    kernel: &'a Vec<u32>,
-    kernel_args: *const u64,
-    launch_bounds: [u32; 3],
-    wave_state: HashMap<usize, WaveState>,
-}
-
-#[derive(Debug, Clone)]
-struct WaveState {
-    scalar_reg: [u32; SGPR_COUNT],
-    scc: u32,
-    vcc: WaveValue,
-    exec: WaveValue,
-    vec_reg: VGPR,
-    pc: usize,
-    sds: HashMap<usize, VecDataStore>,
-}
-
-const SYNCS: [u32; 4] = [0xBF89FC07, 0xBC7C0000, 0xBF890007, 0xbFB60003];
-const S_BARRIER: u32 = 0xBFBD0000;
-
-/// Context for single-stepping through a wave - holds all mutable state
-pub struct WaveContext {
-    pub kernel: Vec<u32>,
-    pub scalar_reg: [u32; SGPR_COUNT],
-    pub scc: u32,
-    pub pc: usize,
-    pub vec_reg: VGPR,
-    pub vcc: WaveValue,
-    pub exec: WaveValue,
-    pub lds: VecDataStore,
-    pub sds: HashMap<usize, VecDataStore>,
-    pub n_lanes: usize,
-}
-
-impl WaveContext {
-    pub fn new(kernel: Vec<u32>, n_lanes: usize) -> Self {
-        let active = (!0u32).wrapping_shr(32 - (n_lanes as u32));
-        Self {
-            kernel,
-            scalar_reg: [0; SGPR_COUNT],
-            scc: 0,
-            pc: 0,
-            vec_reg: VGPR::new(),
-            vcc: WaveValue::new(0, n_lanes),
-            exec: WaveValue::new(active, n_lanes),
-            lds: VecDataStore::new(),
-            sds: (0..=31).map(|i| (i, VecDataStore::new())).collect(),
-            n_lanes,
-        }
-    }
-
-    /// Execute a single instruction. Returns: 0=continue, -1=endpgm, -2=barrier, 1=done (pc past program), negative=error
-    pub fn step(&mut self) -> i32 {
-        if self.pc >= self.kernel.len() { return 1; }
-        if self.kernel[self.pc] == END_PRG { return -1; }
-        if self.kernel[self.pc] == S_BARRIER { self.pc += 1; return -2; }
-        // Skip sync/nop instructions
-        if SYNCS.contains(&self.kernel[self.pc]) || self.kernel[self.pc] >> 20 == 0xbf8 || self.kernel[self.pc] == 0x7E000000 {
-            self.pc += 1;
-            return 0;
-        }
-
-        let mut sgpr_co = None;
-        for lane_id in 0..self.n_lanes {
-            self.vec_reg.default_lane = Some(lane_id);
-            self.vcc.default_lane = Some(lane_id);
-            self.exec.default_lane = Some(lane_id);
-            let mut thread = Thread {
-                scalar_reg: &mut self.scalar_reg,
-                scc: &mut self.scc,
-                vec_reg: &mut self.vec_reg,
-                vcc: &mut self.vcc,
-                exec: &mut self.exec,
-                lds: &mut self.lds,
-                sds: &mut self.sds.get_mut(&lane_id).unwrap(),
-                pc_offset: 0,
-                stream: self.kernel[self.pc..].to_vec(),
-                scalar: false,
-                simm: None,
-                warp_size: self.n_lanes,
-                sgpr_co: &mut sgpr_co,
-            };
-            if let Err(e) = thread.interpret() { return e; }
-            if thread.scalar {
-                self.pc = ((self.pc as isize) + 1 + (thread.pc_offset as isize)) as usize;
-                break;
-            }
-            if lane_id == self.n_lanes - 1 {
-                self.pc = ((self.pc as isize) + 1 + (thread.pc_offset as isize)) as usize;
-            }
-        }
-        if self.vcc.mutations.is_some() { self.vcc.apply_muts(); self.vcc.mutations = None; }
-        if self.exec.mutations.is_some() { self.exec.apply_muts(); self.exec.mutations = None; }
-        if let Some((idx, mut wv)) = sgpr_co.take() { wv.apply_muts(); self.scalar_reg[idx] = wv.value; }
-        0
-    }
-
-    pub fn get_snapshot(&self) -> StateSnapshot {
-        let mut snap = StateSnapshot::new();
-        snap.pc = self.pc as u32;
-        snap.scc = self.scc;
-        snap.vcc = self.vcc.value;
-        snap.exec_mask = self.exec.value;
-        snap.sgpr = self.scalar_reg;
-        for lane in 0..32 { snap.vgpr[lane] = self.vec_reg.get_lane(lane); }
-        snap
-    }
-}
-
-impl<'a> WorkGroup<'a> {
-    pub fn new(dispatch_dim: u32, id: [u32; 3], launch_bounds: [u32; 3], kernel: &'a Vec<u32>, kernel_args: *const u64) -> Self {
-        Self { dispatch_dim, id, kernel, launch_bounds, kernel_args, lds: VecDataStore::new(), wave_state: HashMap::new() }
-    }
-
-    pub fn exec_waves(&mut self) -> Result<(), i32> {
-        let mut threads = vec![];
-        for z in 0..self.launch_bounds[2] {
-            for y in 0..self.launch_bounds[1] {
-                for x in 0..self.launch_bounds[0] {
-                    threads.push([x, y, z])
-                }
-            }
-        }
-        let waves = threads.chunks(WAVE_SIZE).collect::<Vec<_>>();
-
-        let mut sync = false;
-        for (i, x) in self.kernel.iter().enumerate() {
-            if i != 0 && *x == S_BARRIER {
-                sync = true;
-                break;
-            }
-        }
-
-        for _ in 0..=(sync as usize) {
-            for w in waves.iter().enumerate() {
-                self.exec_wave(w)?
-            }
-        }
-        Ok(())
-    }
-
-    fn exec_wave(&mut self, (wave_id, threads): (usize, &&[[u32; 3]])) -> Result<(), i32> {
-        let (mut scalar_reg, mut scc, mut pc, mut vec_reg, mut vcc, mut exec, mut sds) = match self.wave_state.get(&wave_id) {
-          None => {
-            let mut scalar_reg = [0; SGPR_COUNT];
-            scalar_reg.write64(0, self.kernel_args as u64);
-
-            let [gx, gy, gz] = self.id;
-            match self.dispatch_dim {
-              3 => (scalar_reg[13], scalar_reg[14], scalar_reg[15]) = (gx, gy, gz),
-              2 => (scalar_reg[14], scalar_reg[15]) = (gx, gy),
-              _ => scalar_reg[15] = gx,
-            }
-
-            let mut vec_reg = VGPR::new();
-            for (t, [x, y, z]) in threads.iter().enumerate() {
-              vec_reg.get_lane_mut(t)[0] = match &self.launch_bounds {
-                [_, 1, 1] => *x,
-                _ => (z << 20) | (y << 10) | x,
-              }
-            }
-
-            let vcc = WaveValue::new(0, threads.len());
-            let active = (!0u32).wrapping_shr(32 - (threads.len() as u32));
-            let exec = WaveValue::new(active, threads.len());
-
-            let sds = (0..=31).map(|i| (i, VecDataStore::new())).collect();
-            (scalar_reg, 0, 0, vec_reg, vcc, exec, sds)
-          }
-
-          Some(val) => {
-            let val = val.clone();
-            (val.scalar_reg, val.scc, val.pc, val.vec_reg, val.vcc, val.exec, val.sds)
-          }
-        };
-
-        loop {
-            if self.kernel[pc] == END_PRG {
-                break Ok(());
-            }
-            if self.kernel[pc] == S_BARRIER && self.wave_state.get(&wave_id).is_none() {
-                self.wave_state.insert(wave_id, WaveState { scalar_reg, scc, vec_reg, vcc, exec, pc, sds });
-                break Ok(());
-            }
-            if self.kernel[pc] == S_BARRIER || SYNCS.contains(&self.kernel[pc]) || self.kernel[pc] >> 20 == 0xbf8 || self.kernel[pc] == 0x7E000000 {
-                pc += 1;
-                continue;
-            }
-
-            let mut sgpr_co = None;
-            for (lane_id, [x, y, z]) in threads.iter().enumerate() {
-                vec_reg.default_lane = Some(lane_id);
-                vcc.default_lane = Some(lane_id);
-                exec.default_lane = Some(lane_id);
-                if *DEBUG {
-                    let lane = format!("{:<2} {:08X} ", lane_id, self.kernel[pc]);
-                    let state = match exec.read() {
-                        true => "green",
-                        false => "gray",
-                    };
-                    let [id0, id1, id2] = self.id;
-                    print!("[{id0:<3} {id1:<3} {id2:<3}] [{x:<3} {y:<3} {z:<3}] {}", colored(&lane, state));
-                }
-                let mut thread = Thread {
-                    scalar_reg: &mut scalar_reg,
-                    scc: &mut scc,
-                    vec_reg: &mut vec_reg,
-                    vcc: &mut vcc,
-                    exec: &mut exec,
-                    lds: &mut self.lds,
-                    sds: &mut sds.get_mut(&lane_id).unwrap(),
-                    pc_offset: 0,
-                    stream: self.kernel[pc..self.kernel.len()].to_vec(),
-                    scalar: false,
-                    simm: None,
-                    warp_size: threads.len(),
-                    sgpr_co: &mut sgpr_co,
-                };
-                thread.interpret()?;
-                if *DEBUG {
-                    println!();
-                }
-                if thread.scalar {
-                    pc = ((pc as isize) + 1 + (thread.pc_offset as isize)) as usize;
-                    break;
-                }
-                if lane_id == threads.len() - 1 {
-                    pc = ((pc as isize) + 1 + (thread.pc_offset as isize)) as usize;
-                }
-            }
-
-            if vcc.mutations.is_some() {
-                vcc.apply_muts();
-                vcc.mutations = None;
-            }
-            if exec.mutations.is_some() {
-                exec.apply_muts();
-                exec.mutations = None;
-            }
-            if let Some((idx, mut wv)) = sgpr_co.take() {
-                wv.apply_muts();
-                scalar_reg[idx] = wv.value;
-            }
-        }
-    }
-}
-
-#[cfg(test)]
-mod test_workgroup {
-    use super::*;
-
-    // TODO: make this generic by adding the assembler
-    fn global_store_sgpr(addr: u64, instructions: Vec<u32>, src: u32) -> Vec<u32> {
-        [
-            instructions,
-            vec![
-                0x7E020200 + src,
-                0x7E0402FF,
-                addr as u32,
-                0x7E0602FF,
-                (addr >> 32) as u32,
-                0xDC6A0000,
-                0x007C0102,
-            ],
-            vec![END_PRG],
-        ]
-        .concat()
-    }
-
-    #[test]
-    fn test_wave_value_state_vcc() {
-        let mut ret: u32 = 0;
-        let kernel = vec![
-            0xBEEA00FF,
-            0b11111111111111111111111111111111, // initial vcc state
-            0x7E140282,
-            0x7C94010A, // cmp blockDim.x == 2
-        ];
-        let addr = (&mut ret as *mut u32) as u64;
-        let kernel = global_store_sgpr(addr, kernel, 106);
-        let mut wg = WorkGroup::new(1, [0, 0, 0], [3, 1, 1], &kernel, [addr].as_ptr());
-        wg.exec_waves().unwrap();
-        assert_eq!(ret, 0b100);
-    }
-
-    #[test]
-    fn test_wave_value_state_exec() {
-        let mut ret: u32 = 0;
-        let kernel = vec![
-            0xBEFE00FF,
-            0b11111111111111111111111111111111,
-            0x7E140282,
-            0x7D9C010A, // cmpx blockDim.x <= 2
-        ];
-        let addr = (&mut ret as *mut u32) as u64;
-        let kernel = global_store_sgpr(addr, kernel, 126);
-        let mut wg = WorkGroup::new(1, [0, 0, 0], [4, 1, 1], &kernel, [addr].as_ptr());
-        wg.exec_waves().unwrap();
-        assert_eq!(ret, 0b0111);
-    }
-
-    #[test]
-    fn test_wave_value_sgpr_co() {
-        let mut ret: u32 = 0;
-        let kernel = vec![0xBE8D00FF, 0x7FFFFFFF, 0x7E1402FF, u32::MAX, 0xD700000A, 0x0002010A];
-        let addr = (&mut ret as *mut u32) as u64;
-        let kernel = global_store_sgpr(addr, kernel, 0);
-        let mut wg = WorkGroup::new(1, [0, 0, 0], [5, 1, 1], &kernel, [addr].as_ptr());
-        wg.exec_waves().unwrap();
-        assert_eq!(ret, 0b11110);
-    }
-}
@@ -1,155 +0,0 @@
-# ruff: noqa: F405, F403
-# allow define from star imports
-
-import numpy as np
-import unittest
-import subprocess, struct, math, functools
-from tinygrad import Tensor, dtypes, Device
-from tinygrad.helpers import getenv
-
-from tinygrad.runtime.autogen.amd.rdna3.ins import *
-from tinygrad.renderer.amd.asm import waitcnt
-
-from test.testextra.test_cfg_viz import asm_kernel
-
-def get_output(asm:list, n_threads:int=1, vdst:VGPR=v[1]):
-  out = Tensor([0]*n_threads, dtype=dtypes.uint32).realize()
-  insts = [
-    s_load_b64(s[0:1], s[0:1], NULL),
-    *asm,
-    v_lshlrev_b32_e32(v[0], 2, v[0]),
-    s_waitcnt(simm16=waitcnt(lgkmcnt=0)),
-    #global_store_b32(v[0], v[1], s[0:1]),
-    global_store_b32(addr=v[0], data=vdst, saddr=s[0:1]),
-    s_endpgm()
-  ]
-  out = Tensor.custom_kernel(out, fxn=functools.partial(asm_kernel, name="test", insts=insts, device=out.device, n_threads=n_threads))[0]
-  out.realize()
-  return out.tolist()
-
-def f16_to_bits(x:float) -> int: return struct.unpack('<H', struct.pack('<e', x))[0]
-def f32_from_bits(x:int) -> float: return struct.unpack('<f', struct.pack('<I', x))[0]
-def f32_to_bits(x:float) -> int: return struct.unpack('<I', struct.pack('<f', x))[0]
-
-@unittest.skipUnless(Device.DEFAULT == "AMD", "tests RDNA3")
-class TestHW(unittest.TestCase):
-  def setUp(self):
-    if getenv("MOCKGPU"): subprocess.run(["cargo", "build", "--release", "--manifest-path", "./extra/remu/Cargo.toml"], check=True)
-
-  def test_simple_v_mov(self):
-    out = get_output([
-      v_mov_b32_e32(v[1], 2),
-    ])
-    self.assertEqual(out, [2])
-
-  def test_simple_s_mov(self):
-    out = get_output([
-      s_mov_b32(s[7], 0x7fffffff),
-      v_mov_b32_e32(v[1], s[7]),
-    ])
-    self.assertEqual(out, [0x7fffffff])
-
-  def test_exec_mov(self):
-    out = get_output([
-      v_mov_b32_e32(v[1], 42),
-      s_mov_b32(EXEC_LO, 0b10),
-      v_mov_b32_e32(v[1], 10),
-      s_mov_b32(EXEC_LO, 0b11),
-    ], n_threads=2)
-    np.testing.assert_equal(out, [42, 10])
-
-  def test_exec_cmp_vopc(self):
-    out = get_output([
-      s_mov_b32(VCC_LO, 0), # reset vcc
-      v_mov_b32_e32(v[1], 42),
-      v_mov_b32_e32(v[2], 10),
-      s_mov_b32(EXEC_LO, 0b01),
-      v_cmp_ne_u32_e32(v[1], v[2]),
-      s_mov_b32(EXEC_LO, 0b11),
-      v_mov_b32_e32(v[1], VCC_LO),
-    ], n_threads=2)[0]
-    np.testing.assert_equal(out, 1)
-
-  def test_exec_cmpx_vop3(self):
-    out = get_output([
-      s_mov_b32(EXEC_LO, 0b11),
-      v_mov_b32_e32(v[1], 42),
-      v_mov_b32_e32(v[2], 10),
-      s_mov_b32(EXEC_LO, 0b01),
-      v_cmpx_ne_u32_e32(v[1], v[2]),
-      s_mov_b32(s[10], EXEC_LO),
-      s_mov_b32(EXEC_LO, 0b11),
-      v_mov_b32_e32(v[1], s[10]),
-    ], n_threads=2)[0]
-    np.testing.assert_equal(out & 0b11, 0b01)
-
-  def test_fmac_vop3_modifier(self):
-    init_state = [
-      v_mov_b32_e32(a:=v[1], f16_to_bits(4.0)),
-      v_mov_b32_e32(b:=v[2], f16_to_bits(3.0)),
-      v_mov_b32_e32(c:=v[3], f16_to_bits(2.0)),
-    ]
-    def run_fmac(a, b): return get_output(init_state+[v_fmac_f16_e64(c, a, b)], vdst=c)[0]
-    self.assertEqual(run_fmac(a, b), f16_to_bits(14.0))
-    self.assertEqual(run_fmac(a, -b), f16_to_bits(-10.0))
-    self.assertEqual(run_fmac(-a, -b), f16_to_bits(14.0))
-
-  def test_s_abs_i32(self):
-    def check(x, y, dst=s[10], scc=0):
-      for reg,val in [(dst, y), (SCC, scc)]:
-        self.assertEqual(get_output([
-          s_mov_b32(dst, x),
-          s_abs_i32(dst, dst),
-          v_mov_b32_e32(v[1], reg)
-        ])[0], val)
-
-    check(0x00000001, 0x00000001, scc=1)
-    check(0x7fffffff, 0x7fffffff, scc=1)
-    check(0x80000000, 0x80000000, scc=1)
-    check(0x80000001, 0x7fffffff, scc=1)
-    check(0x80000002, 0x7ffffffe, scc=1)
-    check(0xffffffff, 0x00000001, scc=1)
-    check(0, 0, scc=0)
-
-  def test_v_rcp_f32_neg_vop3(self):
-    def v_neg_rcp_f32(x:float, y:float):
-      out = get_output([
-        v_mov_b32_e32(v[2], f32_to_bits(x)),
-        v_rcp_f32_e64(v[2], -v[2]),
-      ], vdst=v[2])[0]
-      assert out == f32_to_bits(y), f"{f32_from_bits(out)} != {y} / {out} != {f32_to_bits(y)}"
-
-    v_neg_rcp_f32(math.inf, -0.0)
-    v_neg_rcp_f32(-math.inf, 0.0)
-    v_neg_rcp_f32(0.0, -math.inf)
-    v_neg_rcp_f32(-0.0, math.inf)
-    v_neg_rcp_f32(-2.0, 0.5)
-    v_neg_rcp_f32(2.0, -0.5)
-
-  def test_v_cndmask_b32_neg(self):
-    def v_neg(x:float, y:float):
-      out = get_output([
-        v_mov_b32_e32(v[1], f32_to_bits(x)),
-        s_mov_b32(s[10], 1),
-        v_cndmask_b32_e64(v[1], v[1], -v[1], s[10]),
-      ])[0]
-      assert out == f32_to_bits(y), f"{f32_from_bits(out)} != {y} / {out} != {f32_to_bits(y)}"
-
-    v_neg(-0.0, 0.0)
-    v_neg(0.0, -0.0)
-    v_neg(2.0, -2.0)
-    v_neg(math.inf, -math.inf)
-    v_neg(-math.inf, math.inf)
-
-  @unittest.skip("how does VOPD work in the dsl")
-  def test_v_subrev_wrap(self):
-    out = get_output([
-      #v_dual_mov_b32(v[1], 0xffffffff, v[2], 0x0),
-      #v_dual_mov_b32(vdstx=v[1], srcx=0xffffffff, vdsty=v[2], srcy=0x0),
-      #VOPD(opx=VOPDOp.V_DUAL_MOV_B32, opy=VOPDOp.V_DUAL_MOV_B32, vdstx=v[1], srcx=0xffffffff, vdsty=v[2], srcy=0x0),
-      v_subrev_co_u32(v[2], VCC_LO, v[2], v[1]),
-    ], vdst=v[2])[0]
-    self.assertEqual(out, 0xffff_ffff)
-
-if __name__ == "__main__":
-  unittest.main()
@@ -0,0 +1,17 @@
+#!/bin/bash
+INSTALL_PATH="${1:-/opt/homebrew/lib}"
+if [ ! -d "$INSTALL_PATH" ]; then
+    USER=$(whoami)
+    echo "No path $INSTALL_PATH. Will create. Might need your password..."
+    echo "You can stop now and provide any location as an argument where you want to save the libs (note, that not default locations should be in LD_LIBRARY_PATH, so tinygrad can find the libs)."
+    echo "Press any key or symbol to continue..."
+    read -n 1 -s
+
+    sudo mkdir -p "$INSTALL_PATH"
+    sudo chown -R "$USER":staff "$INSTALL_PATH"
+fi
+
+# Download libamd_comgr.dylib
+curl -s https://api.github.com/repos/tinygrad/amdcomgr_dylib/releases/latest | \
+    jq -r '.assets[] | select(.name == "libamd_comgr.dylib").browser_download_url' | \
+    xargs curl -L -o $INSTALL_PATH/libamd_comgr.dylib
@@ -3,7 +3,7 @@ INSTALL_PATH="${1:-/opt/homebrew/lib}"
 if [ ! -d "$INSTALL_PATH" ]; then
    USER=$(whoami)
    echo "No path $INSTALL_PATH. Will create. Might need your password..."
-    echo "You can stop now and provide any location as an argument where you want to save the libs (note, that not default locations should be in LD_LIBRARY_PATH, so tinygrad can find the libs)."
+    echo "You can stop now and provide any location as an argument where you want to save the library (note, that not default locations should be in LD_LIBRARY_PATH, so tinygrad can find it)."
    echo "Press any key or symbol to continue..."
    read -n 1 -s

@@ -11,11 +11,6 @@ if [ ! -d "$INSTALL_PATH" ]; then
    sudo chown -R "$USER":staff "$INSTALL_PATH"
 fi

-# Download libremu.dylib
-curl -s https://api.github.com/repos/Qazalin/remu/releases/latest | \
-    jq -r '.assets[] | select(.name == "libremu.dylib").browser_download_url' | \
-    xargs curl -L -o $INSTALL_PATH/libremu.dylib
-
 # Download libamd_comgr.dylib
 curl -s https://api.github.com/repos/tinygrad/amdcomgr_dylib/releases/latest | \
    jq -r '.assets[] | select(.name == "libamd_comgr.dylib").browser_download_url' | \
@@ -0,0 +1,10 @@
+#!/bin/sh
+python3 -c "
+try:
+    from tinygrad.runtime.support.system import APLRemotePCIDevice
+    APLRemotePCIDevice.ensure_app()
+except Exception as e:
+    print('Your tinygrad is too old. Please clone the latest tinygrad: git clone https://github.com/tinygrad/tinygrad.git && cd tinygrad && python3 -m pip install -e .')
+    print(e)
+    exit(1)
+"
@@ -1,6 +1,6 @@
 import os, subprocess, sys, shlex
 from pathlib import Path
-from tinygrad.helpers import temp
+from tinygrad.helpers import temp, getenv

 EXAMPLES_DIR = Path(__file__).parent
 PROFILE_PATH = Path(temp("profile.pkl", append_user=True))
@@ -10,6 +10,7 @@ EXAMPLES = {
  "plus":"test/test_tiny.py TestTiny.test_plus",
  "gemm":"-c \"from tinygrad import Tensor; (Tensor.empty(N:=32, N)@Tensor.empty(N, N)).realize()\"",
  "sync":"test/amd/test_custom_kernel.py TestCustomKernel.test_lds_sync",
+  "handwritten":"test/amd/test_custom_kernel.py TestCustomKernel.test_handwritten",
 }

 if __name__ == "__main__":
@@ -17,9 +18,10 @@ if __name__ == "__main__":
                                 env={**os.environ, "DEBUG":"0"}).rstrip()
  (EXAMPLES_DIR/arch).mkdir(exist_ok=True)
  for name,test in EXAMPLES.items():
+    if getenv("NAME", name) != name: continue
    for i in range(2):
      # AM_RESET=1 gets a clear trace, does not work on mi300 machines
      subprocess.run([sys.executable, *shlex.split(test)], cwd=EXAMPLES_DIR.parent.parent.parent,
-                     env={**os.environ, "AMD":"1", "AM_RESET":"1" if not arch.startswith("gfx9") else "0", "VIZ":"-2", "PYTHONPATH":"."})
+                     env={**os.environ, "DEV":"AMD", "AM_RESET":"1" if not arch.startswith("gfx9") else "0", "VIZ":"-2", "PYTHONPATH":"."})
      PROFILE_PATH.rename(dest:=EXAMPLES_DIR/arch/f"profile_{name}_run_{i}.pkl")
      print(f"saved SQTT trace to {dest}")
--- a/Show More
+++ b/Show More