Merge branch 'master' into algebraic_upat

more mi350x matmul work (#13138 )
* more mi350x matmul work * broken compute
2026-08-19 18:18:27 +00:00 · 2025-11-13 09:21:05 -08:00 · 2025-11-13 09:09:28 -08:00 · 2025-11-13 09:07:54 -08:00 · 2025-11-13 09:07:11 -08:00 · 2025-11-13 08:52:27 -08:00
68 changed files with 5507 additions and 6941 deletions
@@ -61,7 +61,7 @@ runs:
      uses: actions/cache@v4
      with:
        path: ${{ github.workspace }}/.venv
-        key: venv-${{ runner.os }}-python-${{ steps.setup-python.outputs.python-version }}-${{ inputs.deps }}-${{ inputs.pydeps }}-${{ hashFiles('**/setup.py') }}-${{ env.PYTHON_CACHE_VERSION }}
+        key: venv-${{ runner.os }}-python-${{ steps.setup-python.outputs.python-version }}-${{ inputs.deps }}-${{ inputs.pydeps }}-${{ hashFiles('**/pyproject.toml') }}-${{ env.CACHE_VERSION }}

    # **** Caching downloads ****

@@ -70,13 +70,13 @@ runs:
      uses: actions/cache@v4
      with:
        path: ~/.cache/tinygrad/downloads/
-        key: downloads-cache-${{ inputs.key }}-${{ env.DOWNLOAD_CACHE_VERSION }}
+        key: downloads-cache-${{ inputs.key }}-${{ env.CACHE_VERSION }}
    - name: Cache downloads (macOS)
      if: inputs.key != '' && runner.os == 'macOS'
      uses: actions/cache@v4
      with:
        path: ~/Library/Caches/tinygrad/downloads/
-        key: osx-downloads-cache-${{ inputs.key }}-${{ env.DOWNLOAD_CACHE_VERSION }}
+        key: osx-downloads-cache-${{ inputs.key }}-${{ env.CACHE_VERSION }}

    # **** Python deps ****

@@ -187,7 +187,7 @@ runs:
      uses: actions/cache@v4
      with:
        path: /var/cache/apt/archives/
-        key: ${{ runner.os }}-apt-${{ steps.apt-pkgs.outputs.hash }}-${{ env.APT_CACHE_VERSION }}
+        key: ${{ runner.os }}-apt-${{ steps.apt-pkgs.outputs.hash }}-${{ env.CACHE_VERSION }}

    - name: Run apt Update + Install
      if: runner.os == 'Linux' && (inputs.opencl == 'true' || inputs.amd == 'true' || inputs.cuda == 'true' || inputs.webgpu == 'true' || inputs.llvm == 'true')
@@ -247,7 +247,7 @@ runs:
        cache-name: cache-gpuocelot-build-1
      with:
        path: ${{ github.workspace }}/gpuocelot/ocelot
-        key: ${{ runner.os }}-gpuocelot-b16039dc940dc6bc4ea0a98380495769ff35ed99-rebuild-${{ env.BUILD_CACHE_VERSION }}
+        key: ${{ runner.os }}-gpuocelot-b16039dc940dc6bc4ea0a98380495769ff35ed99-rebuild-${{ env.CACHE_VERSION }}
    - name: Clone/compile gpuocelot
      if: inputs.ocelot == 'true' && steps.cache-build.outputs.cache-hit != 'true'
      shell: bash
@@ -1,10 +1,7 @@
 name: Autogen
 env:
  # increment this when downloads substantially change to avoid the internet
-  DOWNLOAD_CACHE_VERSION: '12'
-  PYTHON_CACHE_VERSION: '4'
-  APT_CACHE_VERSION: '1'
-  BUILD_CACHE_VERSION: '1'
+  CACHE_VERSION: '13'
  CAPTURE_PROCESS_REPLAY: 1
  GH_TOKEN: ${{ secrets.GITHUB_TOKEN }}
  PYTHONPATH: ${{ github.workspace }}
@@ -14,10 +11,10 @@ on:
    branches:
      - master
  pull_request:
-    paths: 
+    paths:
    - 'tinygrad/runtime/autogen/**/*'
  workflow_dispatch:
-    paths: 
+    paths:
    - 'tinygrad/runtime/autogen/**/*'

 jobs:
@@ -71,13 +68,10 @@ jobs:
        diff /tmp/sqtt.py.bak tinygrad/runtime/autogen/sqtt.py
    - name: Verify Linux autogen
      run: |
-        cp tinygrad/runtime/autogen/libc.py /tmp/libc.py.bak
        cp tinygrad/runtime/autogen/io_uring.py /tmp/io_uring.py.bak
        cp tinygrad/runtime/autogen/ib.py /tmp/ib.py.bak
-        ./autogen_stubs.sh libc
        ./autogen_stubs.sh io_uring
        ./autogen_stubs.sh ib
-        diff /tmp/libc.py.bak tinygrad/runtime/autogen/libc.py
        diff /tmp/io_uring.py.bak tinygrad/runtime/autogen/io_uring.py
        diff /tmp/ib.py.bak tinygrad/runtime/autogen/ib.py
    - name: Verify WebGPU autogen
@@ -95,3 +89,21 @@ jobs:
        cp tinygrad/runtime/autogen/mesa.py /tmp/mesa.py.bak
        ./autogen_stubs.sh mesa
        diff /tmp/mesa.py.bak tinygrad/runtime/autogen/mesa.py
+  autogen-ng:
+    name: In-tree Autogen
+    runs-on: ubuntu-24.04
+    timeout-minutes: 15
+    steps:
+    - name: Checkout Code
+      uses: actions/checkout@v4
+    - name: Setup Environment
+      uses: ./.github/actions/setup-tinygrad
+      with:
+        pydeps: 'clang>=20'
+    - name: Install autogen support packages
+      run: sudo apt-get install -y --no-install-recommends libclang-20-dev
+    - name: Verify Linux autogen
+      run: |
+        mv tinygrad/runtime/autogen/libc.py /tmp/libc.py.bak
+        python3 -c "from tinygrad.runtime.autogen import libc"
+        diff /tmp/libc.py.bak tinygrad/runtime/autogen/libc.py
@@ -54,7 +54,7 @@ jobs:
    - name: Print macOS version
      run: sw_vers
    - name: Run Stable Diffusion
-      run: BENCHMARK_LOG=stable_diffusion JIT=1 ASSERT_MIN_STEP_TIME=720 python3.11 examples/stable_diffusion.py --fp16 --seed 0 --noshow --timing | tee sd.txt
+      run: BENCHMARK_LOG=stable_diffusion JIT=1 ASSERT_MIN_STEP_TIME=800 python3.11 examples/stable_diffusion.py --fp16 --seed 0 --noshow --timing | tee sd.txt
    - name: Run Stable Diffusion without fp16
      run: BENCHMARK_LOG=stable_diffusion_fp32 JIT=1 ASSERT_MIN_STEP_TIME=800 python3.11 examples/stable_diffusion.py --seed 0 --noshow --timing | tee sd_no_fp16.txt
    - name: Run Stable Diffusion v2
@@ -199,7 +199,7 @@ jobs:
    - name: Test speed vs torch
      run: NV=1 CAPTURE_PROCESS_REPLAY=0 HALF=1 BIG=2 TORCHCUDA=1 python3 test/speed/external_test_speed_v_torch.py | tee torch_speed.txt
    - name: Test speed vs theoretical
-      run: NV=1 IGNORE_BEAM_CACHE=1 DISABLE_COMPILER_CACHE=1 BEAM_DEBUG=1 DEBUG=1 python -m pytest -rA test/external/speed_v_theoretical.py --durations=20
+      run: NV=1 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 benchmark allreduce
      run: NV=1 python test/external/external_benchmark_multitensor_allreduce.py
    - name: Test tensor cores
@@ -320,19 +320,20 @@ jobs:
    #   run: NV=1 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
    - name: Train MNIST
      run: time PYTHONPATH=. NV=1 TARGET_EVAL_ACC_PCT=96.0 python3 examples/beautiful_mnist.py | tee beautiful_mnist.txt
+    # TODO: too slow
    - name: Run 10 CIFAR training steps
-      run: BENCHMARK_LOG=cifar_10steps ASSERT_MIN_STEP_TIME=270 NV=1 STEPS=10 python3 examples/hlb_cifar10.py | tee train_cifar.txt
-    - name: Run 10 CIFAR training steps w HALF
-      run: BENCHMARK_LOG=cifar_10steps_half ASSERT_MIN_STEP_TIME=240 NV=1 STEPS=10 DEFAULT_FLOAT=HALF python3 examples/hlb_cifar10.py | tee train_cifar_half.txt
-    - name: Run 10 CIFAR training steps w BF16
-      run: BENCHMARK_LOG=cifar_10steps_bf16 ASSERT_MIN_STEP_TIME=270 NV=1 STEPS=10 DEFAULT_FLOAT=BFLOAT16 python3 examples/hlb_cifar10.py | tee train_cifar_bf16.txt
+      run: BENCHMARK_LOG=cifar_10steps ASSERT_MIN_STEP_TIME=1300 NV=1 STEPS=10 python3 examples/hlb_cifar10.py | tee train_cifar.txt
+    # - name: Run 10 CIFAR training steps w HALF
+    #   run: BENCHMARK_LOG=cifar_10steps_half ASSERT_MIN_STEP_TIME=240 NV=1 STEPS=10 DEFAULT_FLOAT=HALF python3 examples/hlb_cifar10.py | tee train_cifar_half.txt
+    # - name: Run 10 CIFAR training steps w BF16
+    #   run: BENCHMARK_LOG=cifar_10steps_bf16 ASSERT_MIN_STEP_TIME=270 NV=1 STEPS=10 DEFAULT_FLOAT=BFLOAT16 python3 examples/hlb_cifar10.py | tee train_cifar_bf16.txt
    # TODO: too slow
    # - name: Run 10 CIFAR training steps w winograd
    #   run: BENCHMARK_LOG=cifar_10steps_half_wino ASSERT_MIN_STEP_TIME=350 NV=1 CAPTURE_PROCESS_REPLAY=0 WINO=1 STEPS=10 DEFAULT_FLOAT=HALF python3 examples/hlb_cifar10.py | tee train_cifar_wino.txt
-    - name: Run full CIFAR training w 1 GPU
-      run: time BENCHMARK_LOG=cifar NV=1 DEFAULT_FLOAT=HALF STEPS=1000 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py | tee train_cifar_one_gpu.txt
-    - name: Run full CIFAR training steps w 6 GPUS
-      run: time BENCHMARK_LOG=cifar_6gpu CAPTURE_PROCESS_REPLAY=0 NV=1 DEFAULT_FLOAT=HALF STEPS=350 BS=1536 GPUS=6 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py | tee train_cifar_six_gpu.txt
+    # - name: Run full CIFAR training w 1 GPU
+    #   run: time BENCHMARK_LOG=cifar NV=1 DEFAULT_FLOAT=HALF STEPS=1000 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py | tee train_cifar_one_gpu.txt
+    # - name: Run full CIFAR training steps w 6 GPUS
+    #   run: time BENCHMARK_LOG=cifar_6gpu CAPTURE_PROCESS_REPLAY=0 NV=1 DEFAULT_FLOAT=HALF STEPS=350 BS=1536 GPUS=6 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py | tee train_cifar_six_gpu.txt
    - name: Run MLPerf resnet eval on training data
      run: time BENCHMARK_LOG=resnet_eval NV=1 MODEL=resnet python3 examples/mlperf/model_eval.py
    #- name: Run 10 MLPerf ResNet50 training steps (1 gpu)
@@ -409,7 +410,7 @@ jobs:
    #    python3 -c "import torch; print(torch.__version__)"
    #    LD_PRELOAD="/opt/rocm/lib/libhsa-runtime64.so" HSA=1 BIG=2 TORCHCUDA=1 python3 test/speed/external_test_speed_v_torch.py | tee torch_speed.txt
    - name: Test speed vs theoretical
-      run: AMD=1 IGNORE_BEAM_CACHE=1 DISABLE_COMPILER_CACHE=1 BEAM_DEBUG=1 DEBUG=1 python -m pytest -rA test/external/speed_v_theoretical.py --durations=20
+      run: AMD=1 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
      run: |
        AMD=1 AMD_LLVM=0 python3 test/opt/test_tensor_cores.py
@@ -524,17 +525,18 @@ jobs:
      run: test/external/process_replay/reset.py
    - name: Train MNIST
      run: time PYTHONPATH=. AMD=1 TARGET_EVAL_ACC_PCT=96.0 python3 examples/beautiful_mnist.py | tee beautiful_mnist.txt
+    # TODO: too slow
    - name: Run 10 CIFAR training steps
-      run: BENCHMARK_LOG=cifar_10steps ASSERT_MIN_STEP_TIME=330 AMD=1 STEPS=10 python3 examples/hlb_cifar10.py | tee train_cifar.txt
-    - name: Run 10 CIFAR training steps w HALF
-      run: BENCHMARK_LOG=cifar_10steps_half ASSERT_MIN_STEP_TIME=390 AMD=1 STEPS=10 DEFAULT_FLOAT=HALF python3 examples/hlb_cifar10.py | tee train_cifar_half.txt
+      run: BENCHMARK_LOG=cifar_10steps ASSERT_MIN_STEP_TIME=2000 AMD=1 STEPS=10 python3 examples/hlb_cifar10.py | tee train_cifar.txt
+    # - name: Run 10 CIFAR training steps w HALF
+    #   run: BENCHMARK_LOG=cifar_10steps_half ASSERT_MIN_STEP_TIME=390 AMD=1 STEPS=10 DEFAULT_FLOAT=HALF python3 examples/hlb_cifar10.py | tee train_cifar_half.txt
    # - name: Run 10 CIFAR training steps w BF16
    #   run: BENCHMARK_LOG=cifar_10steps_bf16 ASSERT_MIN_STEP_TIME=288 AMD=1 STEPS=10 DEFAULT_FLOAT=BFLOAT16 python3 examples/hlb_cifar10.py | tee train_cifar_bf16.txt
    # TODO: too slow
    # - name: Run 10 CIFAR training steps w winograd
    #   run: BENCHMARK_LOG=cifar_10steps_half_wino ASSERT_MIN_STEP_TIME=66 AMD=1 WINO=1 STEPS=10 DEFAULT_FLOAT=HALF python3 examples/hlb_cifar10.py | tee train_cifar_wino.txt
-    - name: Run full CIFAR training w 1 GPU
-      run: time BENCHMARK_LOG=cifar AMD=1 DEFAULT_FLOAT=HALF STEPS=1000 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py | tee train_cifar_one_gpu.txt
+    # - name: Run full CIFAR training w 1 GPU
+    #   run: time BENCHMARK_LOG=cifar AMD=1 DEFAULT_FLOAT=HALF STEPS=1000 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py | tee train_cifar_one_gpu.txt
    #- name: Run full CIFAR training steps w 6 GPUS
    #  run: time BENCHMARK_LOG=cifar_6gpu AMD=1 DEFAULT_FLOAT=HALF STEPS=350 BS=1536 GPUS=6 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py | tee train_cifar_six_gpu.txt
    #- name: Run full CIFAR training steps w 6 GPUS (REMOTE)
@@ -623,24 +625,26 @@ 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: openpilot compile3 0.9.9 driving_vision
-      run: BENCHMARK_LOG=openpilot_0_9_9_vision PYTHONPATH=. NOLOCALS=1 FLOAT16=1 IMAGE=2 QCOM=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.9.9/selfdrive/modeld/models/driving_vision.onnx
-    - name: openpilot compile3 0.9.9 driving_policy
-      run: BENCHMARK_LOG=openpilot_0_9_9_policy PYTHONPATH=. NOLOCALS=1 FLOAT16=1 IMAGE=2 QCOM=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.9.9/selfdrive/modeld/models/driving_policy.onnx
-    - name: openpilot compile3 0.9.9 dmonitoring
-      run: BENCHMARK_LOG=openpilot_0_9_9_dmonitoring PYTHONPATH=. NOLOCALS=1 FLOAT16=1 IMAGE=2 QCOM=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.9.9/selfdrive/modeld/models/dmonitoring_model.onnx
+    # - name: openpilot compile3 0.9.9 driving_vision
+    #   run: BENCHMARK_LOG=openpilot_0_9_9_vision PYTHONPATH=. NOLOCALS=1 FLOAT16=1 IMAGE=2 QCOM=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.9.9/selfdrive/modeld/models/driving_vision.onnx
+    # - name: openpilot compile3 0.9.9 driving_policy
+    #   run: BENCHMARK_LOG=openpilot_0_9_9_policy PYTHONPATH=. NOLOCALS=1 FLOAT16=1 IMAGE=2 QCOM=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.9.9/selfdrive/modeld/models/driving_policy.onnx
+    # - name: openpilot compile3 0.9.9 dmonitoring
+    #   run: BENCHMARK_LOG=openpilot_0_9_9_dmonitoring PYTHONPATH=. NOLOCALS=1 FLOAT16=1 IMAGE=2 QCOM=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.9.9/selfdrive/modeld/models/dmonitoring_model.onnx
    - name: openpilot compile3 0.10.0 driving_policy
      run: BENCHMARK_LOG=openpilot_0_10_0_policy PYTHONPATH="." ASSERT_MIN_STEP_TIME=4 DEV=QCOM FLOAT16=1 IMAGE=2 NOLOCALS=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.10.0/selfdrive/modeld/models/driving_policy.onnx
    - name: openpilot compile3 0.10.0 dmonitoring
-      run: BENCHMARK_LOG=openpilot_0_10_0_dmonitoring PYTHONPATH="." ASSERT_MIN_STEP_TIME=12 DEV=QCOM FLOAT16=1 IMAGE=2 NOLOCALS=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.10.0/selfdrive/modeld/models/dmonitoring_model.onnx
+      run: BENCHMARK_LOG=openpilot_0_10_0_dmonitoring PYTHONPATH="." ASSERT_MIN_STEP_TIME=11 DEV=QCOM FLOAT16=1 IMAGE=2 NOLOCALS=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.10.0/selfdrive/modeld/models/dmonitoring_model.onnx
+    - name: DEBUG=2 openpilot compile3 0.10.1 driving_vision
+      run: PYTHONPATH="." DEBUG=2 DEV=QCOM FLOAT16=1 IMAGE=2 NOLOCALS=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/720392c9a5b986981fdbed1bb8c47a6c5573a50e/selfdrive/modeld/models/driving_vision.onnx
    - name: openpilot compile3 0.10.1 driving_vision
      # TODO: ASSERT_MIN_STEP_TIME=17
-      run: BENCHMARK_LOG=openpilot_0_10_1_vision PYTHONPATH="." ASSERT_MIN_STEP_TIME=21 DEV=QCOM FLOAT16=1 IMAGE=2 NOLOCALS=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/720392c9a5b986981fdbed1bb8c47a6c5573a50e/selfdrive/modeld/models/driving_vision.onnx
+      run: BENCHMARK_LOG=openpilot_0_10_1_vision PYTHONPATH="." ASSERT_MIN_STEP_TIME=18 DEV=QCOM FLOAT16=1 IMAGE=2 NOLOCALS=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/720392c9a5b986981fdbed1bb8c47a6c5573a50e/selfdrive/modeld/models/driving_vision.onnx
    - name: openpilot compile3 0.10.1 driving_policy
      run: BENCHMARK_LOG=openpilot_0_10_1_policy PYTHONPATH="." ASSERT_MIN_STEP_TIME=4 DEV=QCOM FLOAT16=1 IMAGE=2 NOLOCALS=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/720392c9a5b986981fdbed1bb8c47a6c5573a50e/selfdrive/modeld/models/driving_policy.onnx
    - name: openpilot compile3 0.10.1 dmonitoring
      # TODO: ASSERT_MIN_STEP_TIME=10
-      run: BENCHMARK_LOG=openpilot_0_10_1_dmonitoring PYTHONPATH="." ASSERT_MIN_STEP_TIME=12 DEV=QCOM FLOAT16=1 IMAGE=2 NOLOCALS=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/720392c9a5b986981fdbed1bb8c47a6c5573a50e/selfdrive/modeld/models/dmonitoring_model.onnx
+      run: BENCHMARK_LOG=openpilot_0_10_1_dmonitoring PYTHONPATH="." ASSERT_MIN_STEP_TIME=11 DEV=QCOM FLOAT16=1 IMAGE=2 NOLOCALS=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/720392c9a5b986981fdbed1bb8c47a6c5573a50e/selfdrive/modeld/models/dmonitoring_model.onnx
    - name: benchmark MobileNetV2 on DSP
      run: |
        # generate quantized weights
@@ -704,8 +708,9 @@ jobs:
      run: |
        AMD=1 GRAPH_ONE_KERNEL=1 PYTHONPATH=. NSZ=8192 python3 test/speed/external_test_copy_speed.py TestCopySpeed.testCopyDefaulttoCPUJit
        AMD=1 GRAPH_ONE_KERNEL=1 PYTHONPATH=. NSZ=8192 python3 test/speed/external_test_copy_speed.py TestCopySpeed.testCopyCPUtoDefaultJit
-    - name: Run full CIFAR training w 1 GPU
-      run: time BENCHMARK_LOG=cifar AMD=1 DEFAULT_FLOAT=HALF STEPS=1000 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py | tee am_train_cifar_one_gpu.txt
+    # TODO: too slow
+    # - name: Run full CIFAR training w 1 GPU
+    #   run: time BENCHMARK_LOG=cifar AMD=1 DEFAULT_FLOAT=HALF STEPS=1000 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py | tee am_train_cifar_one_gpu.txt
    # TODO: enable
    # - name: Run 10 MLPerf ResNet50 training steps (1 gpu)
    #   run: BENCHMARK_LOG=resnet_10steps AMD=1 MNISTMOCK=1 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=256 GPUS=1 MODEL=resnet python3 examples/mlperf/model_train.py | tee am_train_resnet_one_gpu.txt
@@ -767,8 +772,9 @@ jobs:
        NV=1 GRAPH_ONE_KERNEL=1 PYTHONPATH=. NSZ=8192 python3 test/speed/external_test_copy_speed.py TestCopySpeed.testCopyCPUtoDefaultJit
    - name: Test LLAMA-3
      run: BENCHMARK_LOG=llama3_beam NV=1 JITBEAM=2 IGNORE_BEAM_CACHE=1 python3 examples/llama3.py --size 8B --benchmark --temperature 0 | tee nv_llama3_beam.txt
-    - name: Run full CIFAR training w 1 GPU
-      run: time BENCHMARK_LOG=cifar NV=1 DEFAULT_FLOAT=HALF STEPS=1000 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py | tee nv_train_cifar_one_gpu.txt
+    # TODO: too slow
+    # - name: Run full CIFAR training w 1 GPU
+    #   run: time BENCHMARK_LOG=cifar NV=1 DEFAULT_FLOAT=HALF STEPS=1000 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py | tee nv_train_cifar_one_gpu.txt
    #- name: Run 10 MLPerf ResNet50 training steps (1 gpu)
    #  run: BENCHMARK_LOG=resnet_10steps NV=1 MNISTMOCK=1 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=256 GPUS=1 MODEL=resnet python3 examples/mlperf/model_train.py | tee nv_train_resnet_one_gpu.txt
    - name: Run 10 MLPerf Bert training steps (1 gpu)
@@ -22,13 +22,13 @@ jobs:
    - name: Run SDXL with new search
      # TODO: GCVM_L2_PROTECTION_FAULT_STATUS with llvm19
      run: |
-        BENCHMARK_LOG=search_sdxl PYTHONPATH=. AMD=1 JITBEAM=2 IGNORE_BEAM_CACHE=1 DISABLE_COMPILER_CACHE=1 python examples/sdxl.py --noshow --timing --seed 0
+        BENCHMARK_LOG=search_sdxl PYTHONPATH=. AMD=1 JITBEAM=2 IGNORE_BEAM_CACHE=1 CCACHE=0 python examples/sdxl.py --noshow --timing --seed 0
    - name: Run SDXL with cached search
      run: |
        BENCHMARK_LOG=search_sdxl_cached PYTHONPATH=. AMD=1 JITBEAM=2 python examples/sdxl.py --noshow --timing --seed 0
    - name: Run winograd cifar with new search
      run: |
-        BENCHMARK_LOG=search_wino_cifar WINO=1 DEFAULT_FLOAT=HALF JITBEAM=4 IGNORE_BEAM_CACHE=1 DISABLE_COMPILER_CACHE=1 BS=1024 STEPS=500 python examples/hlb_cifar10.py
+        BENCHMARK_LOG=search_wino_cifar WINO=1 DEFAULT_FLOAT=HALF JITBEAM=4 IGNORE_BEAM_CACHE=1 CCACHE=0 BS=1024 STEPS=500 python examples/hlb_cifar10.py
    - name: Run winograd cifar with cached search
      run: |
        BENCHMARK_LOG=search_wino_cifar_cached WINO=1 DEFAULT_FLOAT=HALF JITBEAM=4 BS=1024 STEPS=500 python examples/hlb_cifar10.py
@@ -20,11 +20,11 @@ jobs:
    - name: Install dependencies
      run: |
        python -m pip install --upgrade pip
-        pip install setuptools wheel twine
+        pip install setuptools wheel build twine
    - name: Build and publish
      env:
        TWINE_USERNAME: ${{ secrets.PYPI_USERNAME }}
        TWINE_PASSWORD: ${{ secrets.PYPI_PASSWORD }}
      run: |
-        python setup.py sdist bdist_wheel
+        python -m build
        twine upload dist/*
@@ -1,10 +1,7 @@
 name: Unit Tests
 env:
  # increment this when downloads substantially change to avoid the internet
-  DOWNLOAD_CACHE_VERSION: '12'
-  PYTHON_CACHE_VERSION: '4'
-  APT_CACHE_VERSION: '1'
-  BUILD_CACHE_VERSION: '1'
+  CACHE_VERSION: '13'
  CAPTURE_PROCESS_REPLAY: 1
  GH_TOKEN: ${{ secrets.GITHUB_TOKEN }}
  PYTHONPATH: ${{ github.workspace }}
@@ -233,7 +230,7 @@ jobs:
        python-version: '3.11'
        deps: linting
    - name: Lint bad-indentation and trailing-whitespace with pylint
-      run: python -m pylint --disable=all -e W0311 -e C0303 --jobs=0 --indent-string='  ' --recursive=y .
+      run: python -m pylint --disable=all -e W0311 -e C0303 --jobs=0 --indent-string='  ' --recursive=y . --ignore-paths='tinygrad/runtime/autogen'
    - name: Lint with ruff
      run: |
        pip3 install --upgrade --force-reinstall ruff==0.11.0
@@ -344,10 +341,11 @@ jobs:
          key: gpu-image
          deps: testing_minimal
          opencl: 'true'
-      - name: Test CL IMAGE=2 ops + training
+      - name: Test CL IMAGE=2 ops
        run: |
          CL=1 IMAGE=2 python -m pytest -n=auto test/test_ops.py --durations=20
-          CL=1 IMAGE=2 python test/models/test_end2end.py TestEnd2End.test_linear_mnist
+        # TODO: training is broken
+        # CL=1 IMAGE=2 python test/models/test_end2end.py TestEnd2End.test_linear_mnist
      - name: Run process replay tests
        uses: ./.github/actions/process-replay

@@ -392,7 +390,7 @@ jobs:
          llvm: 'true'
      - name: Test openpilot model kernel count and gate usage
        run: |
-          ALLOWED_KERNEL_COUNT=123 ALLOWED_READ_IMAGE=1452 ALLOWED_GATED_READ_IMAGE=122 FLOAT16=1 CL=1 IMAGE=2 python examples/openpilot/compile3.py https://gitlab.com/commaai/openpilot-lfs.git/gitlab-lfs/objects/cf6376aa9a090f0da26c280ef69eabf9bbdd51d1faac9ed392919c3db69be916
+          ALLOWED_KERNEL_COUNT=123 ALLOWED_READ_IMAGE=1397 ALLOWED_GATED_READ_IMAGE=94 FLOAT16=1 CL=1 IMAGE=2 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 DEBUGCL=1 CL=1 IMAGE=2 python examples/openpilot/compile3.py https://gitlab.com/commaai/openpilot-lfs.git/gitlab-lfs/objects/cf6376aa9a090f0da26c280ef69eabf9bbdd51d1faac9ed392919c3db69be916
      - name: Test openpilot CL compile fp32 (test correctness)
@@ -21,17 +21,38 @@ tinygrad: For something between [PyTorch](https://github.com/pytorch/pytorch) an

 ---

-Despite tinygrad's size, it is a fully featured deep learning framework.
+tinygrad is an end-to-end deep learning stack:

-Due to its extreme simplicity, it is the easiest framework to add new accelerators to, with support for both inference and training. If XLA is CISC, tinygrad is RISC.
+- **Tensor library** with autograd
+- **IR and compiler** that fuse and lower kernels
+- **JIT + graph execution**
+- **nn / optim / datasets** for real training

-tinygrad is now beta software, we [raised some money](https://geohot.github.io/blog/jekyll/update/2023/05/24/the-tiny-corp-raised-5M.html) to make it good. Someday, we will tape out chips.
+It’s inspired by PyTorch (ergonomics), JAX (functional transforms and IR-based AD), and TVM (scheduling and codegen), but stays intentionally tiny and hackable.

-## Features
+---

-### LLaMA and Stable Diffusion
+## How tinygrad compares

-tinygrad can run [LLaMA](/docs/showcase.md#llama) and [Stable Diffusion](/docs/showcase.md#stable-diffusion)!
+**PyTorch**
+
+- ✅ Similar: eager `Tensor` API, autograd, `optim`, basic datasets and layers.
+- ✅ You can write familiar training loops.
+- 🔁 Unlike PyTorch, the entire compiler and IR are visible and hackable.
+
+**JAX**
+
+- ✅ IR-based autodiff over primitives (like JAXPR + XLA).
+- ✅ Function-level JIT (`TinyJit`) that captures and replays kernels.
+- 🔁 Fewer functional transforms (no full `vmap`/`pmap` yet), but far easier to read.
+
+**TVM**
+
+- ✅ Multiple lowering passes, scheduling, and BEAM search over kernels.
+- ✅ Device “graphs” for batched execution.
+- 🔁 tinygrad also ships the **front-end framework** (tensors, nn, optim), not just the compiler.
+
+---

 ### Laziness

@@ -254,23 +254,6 @@ generate_ib() {
  fixup $BASE/ib.py
 }

-generate_libc() {
-  clang2py -k cdefstum \
-    $(dpkg -L libc6-dev | grep sys/mman.h) \
-    $(dpkg -L libc6-dev | grep sys/syscall.h) \
-    /usr/include/string.h \
-    /usr/include/elf.h \
-    /usr/include/unistd.h \
-    /usr/include/asm-generic/mman-common.h \
-    -o $BASE/libc.py
-
-  sed -i "s\import ctypes\import ctypes, ctypes.util, os\g" $BASE/libc.py
-  sed -i "s\FIXME_STUB\libc\g" $BASE/libc.py
-  sed -i "s\FunctionFactoryStub()\None if (libc_path := ctypes.util.find_library('c')) is None else ctypes.CDLL(libc_path, use_errno=True)\g" $BASE/libc.py
-
-  fixup $BASE/libc.py
-}
-
 generate_llvm() {
  INC="$(llvm-config-14 --includedir)"
  clang2py -k cdefstum \
@@ -554,7 +537,6 @@ elif [ "$1" == "sqtt" ]; then generate_sqtt
 elif [ "$1" == "qcom" ]; then generate_qcom
 elif [ "$1" == "io_uring" ]; then generate_io_uring
 elif [ "$1" == "ib" ]; then generate_ib
-elif [ "$1" == "libc" ]; then generate_libc
 elif [ "$1" == "llvm" ]; then generate_llvm
 elif [ "$1" == "kgsl" ]; then generate_kgsl
 elif [ "$1" == "adreno" ]; then generate_adreno
@@ -563,6 +545,6 @@ elif [ "$1" == "vfio" ]; then generate_vfio
 elif [ "$1" == "webgpu" ]; then generate_webgpu
 elif [ "$1" == "libusb" ]; then generate_libusb
 elif [ "$1" == "mesa" ]; then generate_mesa
-elif [ "$1" == "all" ]; then generate_opencl; generate_hip; generate_comgr; generate_cuda; generate_nvrtc; generate_hsa; generate_kfd; generate_nv; generate_amd; generate_io_uring; generate_libc; generate_am; generate_webgpu; generate_mesa
+elif [ "$1" == "all" ]; then generate_opencl; generate_hip; generate_comgr; generate_cuda; generate_nvrtc; generate_hsa; generate_kfd; generate_nv; generate_amd; generate_io_uring; generate_am; generate_webgpu; generate_mesa
 else echo "usage: $0 <type>"
 fi
@@ -1,5 +1,6 @@
 import os, sys, pickle, time, re
 import numpy as np
+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
@@ -0,0 +1,141 @@
+import os
+import numpy as np
+np.set_printoptions(linewidth=1000000)
+os.environ["AMD_LLVM"] = "0"
+
+from tinygrad import Tensor, Context, dtypes, UOp, GlobalCounters
+from tinygrad.helpers import DEBUG, getenv
+from tinygrad.dtype import AddrSpace
+from tinygrad.uop.ops import sint, AxisType, KernelInfo, Ops
+
+WARP_SIZE = 64
+
+# Reg tile sizes (tensor cores)
+TC_M = 16
+TC_N = 16
+TC_K = 32
+
+N,M,K = 4096,4096,4096
+
+# Threadblock tile sizes (block-level tile of C that a block computes)
+BLOCK_M = 64
+BLOCK_N = 64
+BLOCK_K = 64
+
+WARPGROUP_SIZE = 1
+BLOCK_M = BLOCK_M * WARPGROUP_SIZE
+
+TID_SIZE = WARPGROUP_SIZE*WARP_SIZE
+
+def copy(dest:UOp, src:UOp, rng:int, set=False, upcast=()):
+  assert dest.shape == src.shape
+  rngs = [UOp.range(s, rng+i, AxisType.UPCAST if i in upcast else AxisType.LOOP) for i,s in enumerate(src.shape)]
+  copy = dest[*rngs].store(src[*rngs]).end(*rngs)
+  return dest.after(copy) if set else copy
+
+def compute_on_locals(acc:UOp, Asl:UOp, Bsl:UOp, rng:int, afters:tuple[UOp, ...], warpgroup, warp) -> UOp:
+  K_inner_loop = UOp.range(BLOCK_K//TC_K, rng, AxisType.REDUCE)
+
+  # load from locals into registers
+  Ar = UOp.placeholder((BLOCK_M//TC_M//WARPGROUP_SIZE,), dtypes.half.vec(8), slot=1, addrspace=AddrSpace.REG)
+  Br = UOp.placeholder((BLOCK_N//TC_N,), dtypes.half.vec(8), slot=2, addrspace=AddrSpace.REG)
+
+  M_load_loop = UOp.range(BLOCK_M//TC_M//WARPGROUP_SIZE, rng+10)
+  Asl = Asl.reshape(BLOCK_K//TC_K, TC_K, BLOCK_M//TC_M//WARPGROUP_SIZE, WARPGROUP_SIZE, TC_M)
+  load_rng = UOp.range(8, rng+11, axis_type=AxisType.UPCAST)
+  A_in = Asl[K_inner_loop, (warp//16)*8+load_rng, M_load_loop, warpgroup, warp%16].contract(load_rng)
+  Ar = Ar[M_load_loop].set(A_in, end=M_load_loop)
+
+  N_load_loop = UOp.range(BLOCK_N//TC_N, rng+20)
+  Bsl = Bsl.reshape(BLOCK_K//TC_K, TC_K, BLOCK_N//TC_N, TC_N)
+  load_rng = UOp.range(8, rng+21, axis_type=AxisType.UPCAST)
+  B_in = Bsl[K_inner_loop, (warp//16)*8+load_rng, N_load_loop, warp%16].contract(load_rng)
+  Br = Br[N_load_loop].set(B_in, end=N_load_loop)
+
+  M_inner_loop = UOp.range(BLOCK_M//TC_M//WARPGROUP_SIZE, rng+30)
+  N_inner_loop = UOp.range(BLOCK_N//TC_N, rng+31)
+
+  # load values
+  acc_after = acc.after(*afters, M_inner_loop, N_inner_loop, K_inner_loop)
+  acc_load = acc_after[N_inner_loop, M_inner_loop]
+
+  # do WMMA
+  wmma_arg = ('WMMA_16_16_32_half_float', (16, 16, 32), dtypes.half, dtypes.float, 'AMD', 64, ((), (), ((3, 2), (2, 2))), ())
+  out = UOp(Ops.WMMA, dtypes.float.vec(4), (Ar[M_inner_loop], Br[N_inner_loop], acc_load), arg=wmma_arg)
+
+  # store back the acc
+  acc_store = acc[N_inner_loop, M_inner_loop].store(out)
+  return acc_store.end(M_inner_loop, N_inner_loop, K_inner_loop)
+
+def custom_gemm(C:UOp, A:UOp, B:UOp) -> UOp:
+  gx, gy = UOp.special(M//BLOCK_M, "gidx0"), UOp.special(N//BLOCK_N, "gidx1")
+  K_outer_loop = UOp.range(K//BLOCK_K, 0, AxisType.REDUCE)
+
+  # split out the globals into blocks
+  C = C.src[0].cast(dtypes.float.vec(4).ptr(C.ptrdtype.size)).reshape((M//BLOCK_M, BLOCK_M, N//BLOCK_N, BLOCK_N))
+  A = A.reshape((M//BLOCK_M, BLOCK_M, K//BLOCK_K, BLOCK_K))[gx, :, K_outer_loop, :]
+  B = B.reshape((K//BLOCK_K, BLOCK_K, N//BLOCK_N, BLOCK_N))[K_outer_loop, :, gy, :]
+
+  # ---------------------------
+  # GLOBAL -> LOCAL (As, Bs)
+  # ---------------------------
+  tid = UOp.special(TID_SIZE, "lidx0")
+  warpgroup, warp = tid//WARP_SIZE, tid%WARP_SIZE
+
+  A_view = A.reshape(-1, TID_SIZE, 8)
+  B_view = B.reshape(-1, TID_SIZE, 8)
+
+  # A: read BM x BK tiles (permute on store into locals)
+  As = UOp.placeholder((BLOCK_K, BLOCK_M), dtypes.half, slot=0, addrspace=AddrSpace.LOCAL).shrink_to(BLOCK_K, BLOCK_M)
+  As_view = As.reshape(-1, TID_SIZE, 8)
+
+  Bs = UOp.placeholder((BLOCK_K, BLOCK_N+4), dtypes.half, slot=1, addrspace=AddrSpace.LOCAL).shrink_to(BLOCK_K, BLOCK_N)
+  Bs_view = Bs.reshape(-1, TID_SIZE, 8)
+
+  outer_copy = UOp.range(A_view.shape[0], 100, AxisType.UPCAST)
+  inner_copy = UOp.range(A_view.shape[2], 101, AxisType.UPCAST)
+  As_store = As_view[outer_copy, tid, inner_copy].store(A_view[outer_copy, tid, inner_copy])
+  Bs_store = Bs_view[outer_copy, tid, inner_copy].store(B_view[outer_copy, tid, inner_copy])
+
+  if getenv("NOLOAD"):
+    As_store = As[0,0].store(0)
+    Bs_store = Bs[0,0].store(0)
+
+  # TODO: can we automate barrier?
+  barrier = UOp.barrier(UOp.group(As_store, Bs_store).end(outer_copy, inner_copy))
+
+  if getenv("COMPUTE"):
+    As, Bs = As.after(barrier), Bs.after(barrier)
+
+    acc = UOp.placeholder((BLOCK_N//TC_N, BLOCK_M//TC_M//WARPGROUP_SIZE), dtypes.float.vec(4), 0, AddrSpace.REG)
+
+    sink = compute_on_locals(acc, As, Bs, 200, afters=(barrier,), warpgroup=warpgroup, warp=warp)
+    sink = sink.end(K_outer_loop)
+
+    C_view = C[gx, :, gy, :].reshape(BLOCK_M//TC_M//WARPGROUP_SIZE, WARPGROUP_SIZE, TC_M, BLOCK_N//TC_N, TC_N)[:, warpgroup, warp%16, :, (warp//16)*4]
+    sink = copy(C_view, acc.after(sink), rng=300)
+  else:
+    sink = C.after(barrier.end(K_outer_loop))[0,0,0,0].store(As[0,0]+Bs[0,0])
+
+  return sink.sink(arg=KernelInfo(name="custom_gemm", opts_to_apply=())).simplify()
+
+if __name__ == "__main__":
+  a = Tensor.randn(M, K, dtype=dtypes.half)
+  b = Tensor.randn(K, N, dtype=dtypes.half)
+  c = Tensor.empty(M, N, dtype=dtypes.float)
+  with Context(DEBUG=0): Tensor.realize(a,b)
+
+
+  GlobalCounters.reset()
+  with Context(DEBUG=max(2, DEBUG.value), DEVECTORIZE=2):
+    tst = Tensor.custom_kernel(c, a, b, fxn=custom_gemm)[0]
+    tst.realize()
+  print(f"{(N*M*K*2 / GlobalCounters.time_sum_s)*1e-12:.2f} REAL TFLOPS")
+
+
+  with Context(DEBUG=0):
+    ref = a.dot(b, dtype=dtypes.float)
+    ref.realize()
+    #print(ref.numpy())
+    #print(tst.numpy())
+    assert Tensor.isclose(ref, tst, atol=1e-2).all().item(), "matrix not close"
@@ -51,11 +51,15 @@ def create_report(dev, test, result, stdout, stderr):
  dmesg_output = subprocess.check_output(["sudo", "dmesg", "--ctime", "--color=never"], text=True)
  with open(dmesg_path, "w") as f: f.write(dmesg_output)

+  env_vars = " ".join(f"{k}={v}" for k, v in test.env.items())
+  reproduce_cmd = f"{env_vars} {test.cmd}"
+
  summary_path = os.path.join(report_path, "summary.txt")
  with open(summary_path, "w") as f:
    f.write(f"Test: {test.name()}\n")
    f.write(f"Dev params: {vars(dev)}\n")
    f.write(f"Test params: {vars(test)}\n")
+    f.write(f"Reproduce cmd: {reproduce_cmd}\n")
    f.write(f"Exit Code: {result}\n")

  print(f"Crash report saved to {report_path}")
@@ -2,7 +2,8 @@ import os, pathlib, argparse
 from examples.llama3 import Tokenizer
 from tabulate import tabulate
 from tinygrad import fetch
-from tinygrad.helpers import flatten
+from tinygrad.helpers import flatten, getenv
+from sz import NONCORE_DIRS

 # llama 3 tokenizer
 tokenizer = Tokenizer(fetch("https://huggingface.co/bofenghuang/Meta-Llama-3-8B/resolve/main/original/tokenizer.model").as_posix())
@@ -10,19 +11,15 @@ tokenizer = Tokenizer(fetch("https://huggingface.co/bofenghuang/Meta-Llama-3-8B/
 def read_code(base_path):
  ret = []
  for path, _, files in os.walk(os.path.join(base_path, "tinygrad")):
+    if not getenv("CORE") and any(path.split("./")[1].startswith(x) for x in NONCORE_DIRS): continue
    for name in files:
      if not name.endswith(".py"): continue
      if 'tinygrad/runtime/autogen' in path.replace('\\', '/'): continue
      fullpath = os.path.join(path, name)
      code = pathlib.Path(fullpath).read_text()
-      ret.append(("### " + fullpath.split("tinygrad/", 1)[1], code))
+      ret.append((fullpath.split("tinygrad/", 1)[1], code))
  return ret

-def write_code_to_file(filename, code_list):
-  """Writes the combined code to a specified file."""
-  with open(filename, 'w') as f:
-    f.write('\n'.join(flatten(code_list)))
-
 if __name__ == "__main__":
  parser = argparse.ArgumentParser(description="Analyze and optionally save tinygrad code.")
  parser.add_argument("--output", help="Output file to write the combined code to.")
@@ -32,10 +29,11 @@ if __name__ == "__main__":

  table = []
  for name,code in ret:
-    table.append([name, len(tokenizer.encode(name+"\x00"+code))])
+    table.append([name, len(tokenizer.encode(code))])
  print(tabulate([["name", "llm tokens"]]+sorted(table, key=lambda x: -x[1]), headers="firstrow"))

-  code_str = '\x00'.join(flatten(ret))
+  banner = "#"*40
+  code_str = ''.join([f"{banner}\n# {name}\n{banner}\n\n{code}\n" for name,code in ret])
  print(f"code has {len(code_str)} chars")
  newline_count = code_str.count('\n')
  print(f"code has {newline_count} newlines")
@@ -44,5 +42,5 @@ if __name__ == "__main__":
  print(f"code has {len(encoded)} tokens")

  if args.output:
-    write_code_to_file(args.output, ret)
-    print(f"Combined code written to {args.output}")
+    with open(args.output, 'w') as f: f.write(code_str)
+    print(f"Combined code written to {args.output}")
@@ -28,18 +28,6 @@ def llvm_disasm(arch:str, lib:bytes) -> dict[int, tuple[str, int]]:
    cur_off += instr_sz
  return addr_table

-@dataclasses.dataclass
-class InstInfo:
-  typ:str=""
-  inst:str=""
-  hit:int=0
-  lat:int=0
-  stall:int=0
-  def __str__(self): return f"{self.inst:>20} hits:{self.typ:>6} hits:{self.hit:>6} latency:{self.lat:>6} stall:{self.stall:>6}"
-
-  def on_ev(self, ev):
-    self.hit, self.lat, self.stall = self.hit + 1, self.lat + ev.duration, self.stall + ev.stall
-
@dataclasses.dataclass(frozen=True)
 class InstExec:
  typ:str
@@ -48,25 +36,18 @@ class InstExec:
  dur:int
  time:int

-@dataclasses.dataclass(frozen=True)
-class PrgExec:
-  name:str
-  wave:int
-  cu:int
-  simd:int
-  def __str__(self): return f"{self.name},{self.wave},{self.cu},{self.simd}"
-
@dataclasses.dataclass(frozen=True)
 class WaveExec:
  wave_id:int
  cu:int
  simd:int
+  begin_time:int
+  end_time:int
  insts:list[InstExec]

 class _ROCParseCtx:
  def __init__(self, dev_evs:dict[str, ProfileDeviceEvent], sqtt_evs:list[ProfileSQTTEvent], prog_evs:list[ProfileProgramEvent]):
    self.dev_evs, self.sqtt_evs, self.prog_evs = dev_evs, iter(sqtt_evs), prog_evs
-    self.wave_events:dict[PrgExec, dict[int, InstInfo]] = {}
    self.disasms:dict[tuple[str, int], tuple[str, int]] = {}
    self.inst_execs:dict[str, list[WaveExec]] = {}

@@ -79,7 +60,8 @@ class _ROCParseCtx:
    x = next(self.sqtt_evs, None)
    self.active_kern = x.kern if x is not None else None
    self.active_se = x.se if x is not None else None
-    return x
+    self.active_blob = (ctypes.c_ubyte * len(x.blob)).from_buffer_copy(x.blob) if x is not None else None
+    return self.active_blob

  def on_occupancy_ev(self, ev):
    if DEBUG >= 5: print("OCC", ev.time, self.active_se, ev.cu, ev.simd, ev.wave_id, ev.start)
@@ -87,19 +69,15 @@ class _ROCParseCtx:
  def on_wave_ev(self, ev):
    if DEBUG >= 5: print("WAVE", ev.wave_id, self.active_se, ev.cu, ev.simd, ev.contexts, ev.begin_time, ev.end_time)

-    asm:dict[int, InstInfo] = {}
    inst_execs:list[InstExec] = []
    for j in range(ev.instructions_size):
      inst_ev = ev.instructions_array[j]
      inst_typ = rocprof.rocprofiler_thread_trace_decoder_inst_category_t__enumvalues[inst_ev.category]
      inst_disasm = self.disasms[(unwrap(self.active_kern), unwrap(inst_ev.pc.address))][0]
-      asm.setdefault(inst_ev.pc.address, InstInfo(typ=inst_typ, inst=inst_disasm))
-      asm[inst_ev.pc.address].on_ev(inst_ev)
      inst_execs.append(InstExec(inst_typ, inst_disasm, inst_ev.stall, inst_ev.duration, inst_ev.time))

    if ev.instructions_size > 0:
-      self.wave_events[key:=PrgExec(unwrap(self.active_kern), ev.wave_id, ev.cu, ev.simd)] = asm
-      self.inst_execs.setdefault(key.name, []).append(WaveExec(ev.wave_id, ev.cu, ev.simd, inst_execs))
+      self.inst_execs.setdefault(unwrap(self.active_kern), []).append(WaveExec(ev.wave_id, ev.cu, ev.simd, ev.begin_time, ev.end_time, inst_execs))

 def decode(profile:list[ProfileEvent]) -> _ROCParseCtx:
  dev_events:dict[str, ProfileDeviceEvent] = {}
@@ -114,10 +92,10 @@ def decode(profile:list[ProfileEvent]) -> _ROCParseCtx:

  @rocprof.rocprof_trace_decoder_se_data_callback_t
  def copy_cb(buf, buf_size, data_ptr):
-    if (prof:=ROCParseCtx.next_sqtt()) is None: return 0
-    buf[0] = ctypes.cast((ctypes.c_ubyte * len(prof.blob)).from_buffer_copy(prof.blob), ctypes.POINTER(ctypes.c_ubyte))
-    buf_size[0] = len(prof.blob)
-    return len(prof.blob)
+    if (prof_info:=ROCParseCtx.next_sqtt()) is None: return 0
+    buf[0] = ctypes.cast(prof_info, ctypes.POINTER(ctypes.c_ubyte))
+    buf_size[0] = len(prof_info)
+    return len(prof_info)

  @rocprof.rocprof_trace_decoder_trace_callback_t
  def trace_cb(record_type, events_ptr, n, data_ptr):
@@ -147,7 +125,7 @@ def decode(profile:list[ProfileEvent]) -> _ROCParseCtx:

  try:
    rocprof.rocprof_trace_decoder_parse_data(copy_cb, trace_cb, isa_cb, None)
-  except AttributeError as e: raise RuntimeError("Failed to find rocprof-trace-decoder. Run ./extra/sqtt/install_sqtt_decoder.py to install") from e
+  except AttributeError as e: raise RuntimeError("Failed to find rocprof-trace-decoder. Run sudo ./extra/sqtt/install_sqtt_decoder.py to install") from e
  return ROCParseCtx

 if __name__ == "__main__":
@@ -157,7 +135,7 @@ if __name__ == "__main__":

  with args.profile.open("rb") as f: profile = pickle.load(f)
  rctx = decode(profile)
-  print('SQTT:', rctx.wave_events.keys())
+  print('SQTT:', rctx.inst_execs.keys())

  for ev in profile:
    if not isinstance(ev, ProfilePMCEvent): continue
@@ -1,7 +1,7 @@
 import os
 os.environ["PYTHONPATH"] = "."
 os.environ["SQTT"] = "1"
-os.environ["AMD"] = "1"
+if "DEV" not in os.environ: os.environ["DEV"] = "AMD"
 os.environ["VIZ"] = "1"
 os.environ["AMD_LLVM"] = "0"

@@ -10,13 +10,13 @@ import sys, contextlib
 from tinygrad import Tensor
 from tinygrad.dtype import dtypes
 from tinygrad.renderer import ProgramSpec
-from tinygrad.uop.ops import UOp, Ops, KernelInfo
+from tinygrad.uop.ops import UOp, Ops, KernelInfo, AddrSpace
 from tinygrad.engine.realize import CompiledRunner
 from tinygrad.device import Device, ProfileDeviceEvent

 from extra.sqtt.roc import decode, InstExec, PrgExec

-dev = Device["AMD"]
+dev = Device[os.environ["DEV"]]

 def custom(arg:str, s:UOp|None=None) -> UOp: return UOp(Ops.CUSTOM, src=(s,) if s is not None else (), arg=arg)

@@ -39,9 +39,10 @@ def save_sqtt():
  sqtt:dict[PrgExec, list[InstExec]] = {}
  yield sqtt
  # decode sqtt
-  rctx = decode(dev.profile_events+[ProfileDeviceEvent("AMD", props=dev.device_props())])
-  assert len(rctx.inst_execs) > 0, "empty sqtt output"
-  sqtt.update(rctx.inst_execs)
+  if os.environ["DEV"] == "AMD":
+    rctx = decode(dev.profile_events+[ProfileDeviceEvent("AMD", props=dev.device_props())])
+    assert len(rctx.inst_execs) > 0, "empty sqtt output"
+    sqtt.update(rctx.inst_execs)

 class TestTiming(unittest.TestCase):
  def test_v_add(self):
@@ -62,18 +63,24 @@ class TestTiming(unittest.TestCase):
    assert all(s.stall == 0 for s in wave)

  def test_multi_cycle_inst(self):
+    def custom_vrcp(A, B):
+      op = custom("float a = 0.0;")
+      op = custom("float b = (*(data1_1+0));", op)
+      #op = custom('asm volatile("v_mul_f32_e32 %2 %2 %1" : "+v"(a) : "v"(b));', op)
+      op = custom('asm volatile("v_rcp_f32_e32 %2 %1" : "+v"(a) : "v"(b));', op)
+      op = custom('asm volatile("v_add_f32_e64 %1 %1 1.0" : "+v"(a));', op)
+      op = custom("*(data0_1+0) = a;", op)
+      return UOp.sink(op, A, B, arg=KernelInfo(name="custom_vrcp"))
+    out = Tensor([0.]).realize()
+    inp = Tensor([-2.0]).realize()
    with save_sqtt() as sqtt:
-      asm_kernel([
-        "v_mov_b32_e32 v4 0x3f800000",
-        "v_rcp_f32_e32 v5 v4",
-        "v_mul_f32_e32 v6 v5 v4",
-      ]).realize()
-    w = list(sqtt.values())[0]
-    rcp, mul = w[1], w[2]
-    self.assertGreater(rcp.dur, 1) # 4 cycles on gfx11
-    self.assertEqual(mul.dur, 1)
-    # mul depends on v5, how can it run before rcp is done?
-    self.assertGreaterEqual(mul.time, rcp.time+rcp.dur)
+      Tensor.custom_kernel(out, inp, fxn=custom_vrcp)[0].realize()
+
+    wave = list(sqtt.values())[0][0]
+    for i in range(len(wave.insts)):
+      if wave.insts[i].inst.startswith("global_store"):
+        print(f"store diff {wave.insts[i].time-(wave.insts[i-1].time)}")
+    self.assertEqual(out.item(), 0.5)

  def test_wmma(self):
    with save_sqtt() as sqtt:
@@ -7,7 +7,7 @@ from tinygrad.dtype import AddrSpace, PtrDType
 from tinygrad.helpers import getenv, prod

 from extra.thunder.tiny.tk import WARP_THREADS
-from extra.thunder.tiny.tk.tiles import TILE_ROW_DIM, TILE_COL_DIM, RT_BASE_TILE_NEPT, slots
+from extra.thunder.tiny.tk.tiles import RT

 class Group:
  def __init__(self, warps:int, ker):
@@ -32,7 +32,11 @@ class Group:

    i = UOp.range(reg.size, Group.clear_rid)
    Group.clear_rid += 1
-    return reg.reshape((reg.size,))[i].set(value, end=i).after(reg).reshape(reg.shape)
+
+    reg_store = reg.reshape((reg.size,))[i].store(value).end(i)
+
+    self.ker.push_store(reg_store, reg)
+    return reg.after(reg_store).reshape(reg.shape)

  def zero(self, reg:UOp): return self.clear(reg, 0)
  def neg_inf(self, reg:UOp): return self.clear(reg, -math.inf)
@@ -122,27 +126,38 @@ class Group:
  def row_reduce(self, vec:UOp, src:UOp, op:Callable[[UOp, UOp], UOp]):
    assert self.warps == 1

-    red_local = UOp.placeholder((self.group_threads, 2), src.dtype.base, addrspace=AddrSpace.LOCAL, slot=slots.shared_slot)
-    slots.shared_slot += 1
+    red_local = self.ker.alloc((self.group_threads, 2), src.dtype.base, AddrSpace.LOCAL)
+    red_reg = self.ker.alloc((2,), src.dtype.base, AddrSpace.REG)

    for height in self.ker.range(src.shape[-3], track=False):
+      i = UOp.range(red_reg.size, Group.clear_rid)
+      Group.clear_rid += 1
+      red_reg = red_reg.after(height, *[tkr._rng for tkr in self.ker.range_stack])
+      reg_store = red_reg.flatten()[i].store(0.).end(i)
+      red_reg = red_reg.after(reg_store).reshape(red_reg.shape)
+
      for i_outer in self.ker.range(2, track=False):
        for width in self.ker.range(src.shape[-2], AxisType.REDUCE, track=False):
          for i_inner in self.ker.range(4, AxisType.REDUCE, track=False):
            elem_index = i_inner + 2 * (i_inner // 2) + i_outer * 2
-            vec_store = vec[height, 0, i_outer].store(op(vec[height, 0, i_outer], src[height, width, elem_index])).end(width, i_inner, i_outer)
-            vec = vec.after(vec_store).reshape(vec.shape)
+            reg_store = red_reg[i_outer].store(op(red_reg[i_outer], src[height, width, elem_index])).end(i_inner, width, i_outer)
+            red_reg = red_reg.after(reg_store).reshape(red_reg.shape)

      # store to shared memory
      for i_outer in self.ker.range(2, track=False):
-        red_local_store = red_local[self.laneid, i_outer].store(vec[height, 0, i_outer]).end(i_outer)
-        red_local = red_local.after(red_local_store).reshape(red_local.shape)
+        red_local_store = red_local[self.laneid, i_outer].store(red_reg[i_outer]).end(i_outer)
+        red_local = red_local.after(red_local_store.barrier()).reshape(red_local.shape)

      # reduce from shared memory
      for i_outer in self.ker.range(2, track=False):
        for i_inner in self.ker.range(3, AxisType.REDUCE, track=False):
-          offset = (self.laneid // 4) * 4 + ((self.laneid + 1 + i_inner) % 4)
-          vec_store = vec[height, 0, i_outer].store(op(vec[height, 0, i_outer], red_local[offset, i_outer])).end(i_inner, i_outer)
+          offset = (self.laneid // 4) * 4 + ((self.laneid + i_inner + 1) % 4)
+          reg_store = red_reg[i_outer].store(op(red_reg[i_outer], red_local[offset, i_outer])).end(i_inner, i_outer)
+          red_reg = red_reg.after(reg_store).reshape(red_reg.shape)
+
+      # reduce with vec
+      for i_outer in self.ker.range(2, track=False):
+        vec_store = vec[height, 0, i_outer].store(op(vec[height, 0, i_outer], red_reg[i_outer])).end(i_outer, height)

    self.ker.push_store(vec_store, vec)
    return vec.after(vec_store).reshape(vec.shape)
@@ -159,7 +174,7 @@ class Group:

      load_i_height = UOp.range(dst.shape[-3], Group.load_rid)
      load_i_width = UOp.range(dst.shape[-2], Group.load_rid+1)
-      load_i_inner = UOp.range(RT_BASE_TILE_NEPT, Group.load_rid+2)
+      load_i_inner = UOp.range(RT.BASE_TILE_NEPT, Group.load_rid+2)
      Group.load_rid += 3

      if self.warps % 4 == 0: local_warpid = (self.warpid // 4) + (self.warpid % 4) * (self.warps // 4)
@@ -167,14 +182,14 @@ class Group:
      warp_laneid = self.threadIdx_x % WARP_THREADS

      if not transpose:
-        row = (local_warpid * dst.shape[-3] + load_i_height) * TILE_ROW_DIM + (warp_laneid // 4)
-        col = load_i_width * TILE_COL_DIM + 2 * (warp_laneid % 4)
+        row = (local_warpid * dst.shape[-3] + load_i_height) * RT.TILE_ROW_DIM + (warp_laneid // 4)
+        col = load_i_width * RT.TILE_COL_DIM + 2 * (warp_laneid % 4)

        row_offset = ((load_i_inner % 4) // 2) * 8
        col_offset = (load_i_inner % 2) + (load_i_inner // 4) * 8
      else:
-        row = (local_warpid * dst.shape[-3] + load_i_height) * TILE_ROW_DIM + 2 * (warp_laneid % 4)
-        col = load_i_width * TILE_COL_DIM + (warp_laneid // 4)
+        row = (local_warpid * dst.shape[-3] + load_i_height) * RT.TILE_ROW_DIM + 2 * (warp_laneid % 4)
+        col = load_i_width * RT.TILE_COL_DIM + (warp_laneid // 4)

        row_offset = (load_i_inner % 2) + (load_i_inner // 4) * 8
        col_offset = ((load_i_inner % 4) // 2) * 8
@@ -223,15 +238,15 @@ class Group:

      store_i_height = UOp.range(src.shape[-3], Group.store_rid)
      store_i_width = UOp.range(src.shape[-2], Group.store_rid+1)
-      store_i_inner = UOp.range(RT_BASE_TILE_NEPT, Group.store_rid+2)
+      store_i_inner = UOp.range(RT.BASE_TILE_NEPT, Group.store_rid+2)
      Group.store_rid += 3

      if self.warps % 4 == 0: local_warpid = (self.warpid // 4) + (self.warpid % 4) * (self.warps // 4)
      else: local_warpid = self.warpid
      warp_laneid = self.threadIdx_x % WARP_THREADS

-      row = (local_warpid * src.shape[-3] + store_i_height) * TILE_ROW_DIM + (warp_laneid // 4)
-      col = store_i_width * TILE_COL_DIM + 2 * (warp_laneid % 4)
+      row = (local_warpid * src.shape[-3] + store_i_height) * RT.TILE_ROW_DIM + (warp_laneid // 4)
+      col = store_i_width * RT.TILE_COL_DIM + 2 * (warp_laneid % 4)

      row_offset = ((store_i_inner % 4) // 2) * 8
      col_offset = (store_i_inner % 2) + (store_i_inner // 4) * 8
@@ -1,7 +1,8 @@
 from contextlib import AbstractContextManager
-from tinygrad.uop.ops import UOp, KernelInfo, AxisType
+from tinygrad.uop.ops import UOp, KernelInfo, AxisType, AddrSpace
 from extra.thunder.tiny.tk import WARP_THREADS
 from extra.thunder.tiny.tk.group import Group
+from extra.thunder.tiny.tk.tiles import GL, ST, RT, RV

 class _tk_range:
  user_rid = 0
@@ -25,6 +26,11 @@ class Kernel(AbstractContextManager):
    self.range_stack = []
    self.store_stack = []

+    self.global_slot = 0
+    self.shared_slot = 0
+    self.register_slot = 0
+    self.allocs = {}
+
  @property
  def warpid(self): return self.threadIdx_x // WARP_THREADS

@@ -42,6 +48,31 @@ class Kernel(AbstractContextManager):
    if track: self.range_stack.append(rng)
    return rng

+  def alloc(self, shape, dtype, addrspace:AddrSpace, name:str|None=None):
+    match addrspace:
+      case AddrSpace.GLOBAL:
+        slot = self.global_slot
+        self.global_slot += 1
+      case AddrSpace.LOCAL:
+        slot = self.shared_slot
+        self.shared_slot += 1
+      case AddrSpace.REG:
+        slot = self.register_slot
+        self.register_slot += 1
+
+    uop = UOp.placeholder(shape, dtype, slot=slot, addrspace=addrspace)
+
+    if name:
+      if (name, shape) in self.allocs: return self.allocs[(name, shape)]
+      self.allocs[(name, shape)] = uop
+
+    return uop
+
+  def gl(self, shape, dtype): return GL(shape, dtype, self)._uop
+  def st(self, shape, dtype): return ST(shape, dtype, self)._uop
+  def rt(self, shape, dtype): return RT(shape, dtype, self)._uop
+  def rv(self, length, dtype, layout="naive"): return RV(length, dtype, layout, self)._uop
+
  def push_store(self, store:UOp, uop:UOp): self.store_stack.append((store, uop))

  def finish(self):
@@ -1,52 +1,45 @@
-import math
-from typing import cast, Callable
-from tinygrad import Tensor, Device, Context, GlobalCounters, dtypes
-from tinygrad.uop.ops import AxisType, UOp, KernelInfo, Ops
-from tinygrad.engine.realize import ExecItem, get_runner
-from tinygrad.dtype import AddrSpace, PtrDType
-from tinygrad.helpers import getenv, prod
+from tinygrad.dtype import AddrSpace

 from extra.thunder.tiny.tk import WARP_THREADS

-class _Slots:
-  def __init__(self):
-    self.global_slot = 0
-    self.shared_slot = 0
-    self.register_slot = 0
-slots = _Slots()
+class GL:
+  def __init__(self, shape, dtype, ker):
+    self.shape, self.dtype = shape, dtype
+    self._uop = ker.alloc(shape, dtype, AddrSpace.GLOBAL)

-def gl(shape, dtype):
-  slots.global_slot += 1
-  return UOp.placeholder(shape, dtype, slot=slots.global_slot-1)
+class ST:
+  def __init__(self, shape, dtype, ker):
+    self.shape, self.dtype = shape, dtype
+    self._uop = ker.alloc(shape, dtype, AddrSpace.LOCAL)

-shared_slot = 0
-def st(shape, dtype):
-  slots.shared_slot += 1
-  return UOp.placeholder(shape, dtype, addrspace=AddrSpace.LOCAL, slot=slots.shared_slot-1)
+class RT:
+  TILE_ROW_DIM, TILE_COL_DIM = 16, 16
+  BASE_TILE_NE = TILE_ROW_DIM * TILE_COL_DIM
+  BASE_TILE_NEPT = BASE_TILE_NE // WARP_THREADS

-TILE_ROW_DIM, TILE_COL_DIM = 16, 16
-RT_BASE_TILE_NE = TILE_ROW_DIM * TILE_COL_DIM
-RT_BASE_TILE_NEPT = RT_BASE_TILE_NE // WARP_THREADS
-register_slot = 0
-def rt(shape, dtype):
-  assert len(shape) == 2
+  def __init__(self, shape, dtype, ker):
+    assert len(shape) == 2
+    assert shape[0] % RT.TILE_ROW_DIM == 0
+    assert shape[1] % RT.TILE_COL_DIM == 0

-  height = shape[0] // TILE_ROW_DIM
-  width = shape[1] // TILE_COL_DIM
+    height = shape[0] // RT.TILE_ROW_DIM
+    width = shape[1] // RT.TILE_COL_DIM

-  slots.register_slot += 1
-  return UOp.placeholder((height, width, RT_BASE_TILE_NEPT), dtype, addrspace=AddrSpace.REG, slot=slots.register_slot-1)
+    self.shape, self.dtype = (height, width, self.BASE_TILE_NEPT), dtype
+    self._uop = ker.alloc(self.shape, dtype, AddrSpace.REG)

-def rv(length, dtype, layout="naive"):
-  tiles = length // TILE_ROW_DIM
-  match layout:
-    case "naive":
-      inner_dim = 1
-      outer_dim = (tiles + 1) // 2
-    case "ortho":
-      inner_dim = 1
-      outer_dim = tiles
-    case _: raise NotImplementedError(f"rv layout {layout} not implemented")
+class RV:
+  def __init__(self, length, dtype, layout, ker):
+    tiles = length // RT.TILE_ROW_DIM

-  slots.register_slot += 1
-  return UOp.placeholder((outer_dim, inner_dim, 2), dtype, addrspace=AddrSpace.REG, slot=slots.register_slot-1)
+    match layout:
+      case "naive":
+        inner_dim = 1
+        outer_dim = (tiles + 1) // 2
+      case "ortho":
+        inner_dim = 1
+        outer_dim = tiles
+      case _: raise NotImplementedError(f"rv layout {layout} not implemented")
+
+    self.shape, self.dtype = (outer_dim, inner_dim, 2), dtype
+    self._uop = ker.alloc(self.shape, dtype, AddrSpace.REG)
@@ -1,6 +1,5 @@
 #!/usr/bin/env python3
 import sys, os, zlib, struct, hashlib
-from hexdump import hexdump
 from tinygrad.helpers import DEBUG, getenv, fetch
 from tinygrad.runtime.support.usb import USB3

@@ -1,10 +0,0 @@
-[mypy]
-warn_unused_configs = True
-files = tinygrad
-ignore_missing_imports = True
-check_untyped_defs = True
-explicit_package_bases = True
-warn_unreachable = True
-warn_redundant_casts = True
-# NOTE: had to comment this out to make mypy pass on both CI and OSX
-#warn_unused_ignores = True
@@ -0,0 +1,233 @@
+[project]
+name = "tinygrad"
+version = "0.11.0"
+description = "You like pytorch? You like micrograd? You love tinygrad! <3"
+authors = [{ name = "George Hotz" }]
+
+classifiers = ["Programming Language :: Python :: 3"]
+
+license = 'MIT'
+readme = "README.md"
+requires-python = ">=3.11"
+dependencies = []
+
+[build-system]
+requires = ["setuptools"]
+build-backend = "setuptools.build_meta"
+
+[tool.setuptools]
+include-package-data = true
+packages = [
+  'tinygrad',
+  'tinygrad.apps',
+  'tinygrad.codegen',
+  'tinygrad.codegen.opt',
+  'tinygrad.codegen.late',
+  'tinygrad.engine',
+  'tinygrad.mixin',
+  'tinygrad.nn',
+  'tinygrad.renderer',
+  'tinygrad.runtime',
+  'tinygrad.runtime.autogen',
+  'tinygrad.runtime.autogen.am',
+  'tinygrad.runtime.autogen.nv',
+  'tinygrad.runtime.graph',
+  'tinygrad.runtime.support',
+  'tinygrad.runtime.support.am',
+  'tinygrad.runtime.support.nv',
+  'tinygrad.schedule',
+  'tinygrad.uop',
+  'tinygrad.viz',
+]
+
+[tool.setuptools.package-data]
+tinygrad = ["py.typed"]
+"tinygrad.viz" = ["index.html", "assets/**/*", "js/*"]
+
+
+[project.optional-dependencies]
+arm = ["unicorn"]
+triton = ["triton-nightly>=2.1.0.dev20231014192330"]
+linting = [
+  "pylint",
+  "mypy==1.18.1",
+  "typing-extensions",
+  "pre-commit",
+  "ruff",
+  "numpy",
+  "typeguard",
+]
+# mlperf = [
+#   "mlperf-logging @ git+https://github.com/mlperf/[email protected]",
+# ]
+testing_minimal = [
+  "numpy",
+  "torch==2.9.0",
+  "pytest",
+  "pytest-xdist",
+  "pytest-timeout",
+  "pytest-split",
+  "hypothesis",
+  "z3-solver",
+]
+testing_unit = ["tinygrad[testing_minimal]", "tqdm", "safetensors", "tabulate"]
+testing = [
+  "tinygrad[testing_minimal]",
+  "pillow",
+  "onnx==1.18.0",
+  "onnx2torch",
+  "onnxruntime",
+  "opencv-python",
+  "tabulate",
+  "tqdm",
+  "safetensors",
+  "transformers",
+  "sentencepiece",
+  "tiktoken",
+  "blobfile",
+  "librosa",
+  # librosa needs numba but uv ignores python upper bounds and some numba versions require <python3.10
+  "numba>=0.55",
+  "networkx",
+  "nibabel",
+  "bottle",
+  "ggml-python",
+  "capstone",
+  "pycocotools",
+  "boto3",
+  "pandas",
+  "influxdb3-python",
+]
+docs = [
+  "mkdocs",
+  "mkdocs-material",
+  "mkdocstrings[python]",
+  "markdown-callouts",
+  "markdown-exec[ansi]",
+  "black",
+  "numpy",
+]
+
+
+[tool.mutmut]
+paths_to_mutate = ["tinygrad/"]
+do_not_mutate = [
+  "tinygrad/apps/*",
+  "tinygrad/codegen/*",
+  "tinygrad/engine/*",
+  "tinygrad/nn/*",
+  "tinygrad/renderer/*",
+  "tinygrad/runtime/*",
+  "tinygrad/schedule/*",
+  "tinygrad/uop/*",
+  "tinygrad/viz/*",
+  "tinygrad/device.py",
+  "tinygrad/dtype.py",
+  "tinygrad/gradient.py",
+  "tinygrad/helpers.py",
+  "tinygrad/tensor.py",
+]
+tests_dir = ["test/test_tiny.py", "test/test_ops.py"]
+debug = true
+
+
+[tool.mypy]
+warn_unused_configs = true
+files = ["tinygrad"]
+ignore_missing_imports = true
+check_untyped_defs = true
+explicit_package_bases = true
+warn_unreachable = true
+warn_redundant_casts = true
+# NOTE: had to comment this out to make mypy pass on both CI and OSX
+#warn_unused_ignores = true
+
+[tool.pytest.ini_options]
+norecursedirs = [
+  "extra",
+  ".hypothesis",
+  ".git",
+]
+timeout = 300
+timeout_method = "thread"
+timeout_func_only = true
+testpaths = ["test"]
+
+[tool.ruff]
+preview = true
+target-version = "py311"
+line-length = 150
+indent-width = 2
+exclude = [
+  ".git/",
+  "docs/",
+  "extra/",
+  "tinygrad/runtime/autogen",
+  "test/external/mlperf_resnet",
+  "test/external/mlperf_unet3d",
+]
+
+[tool.ruff.lint]
+select = [
+  "F",  # Pyflakes
+  "W6",
+  "E71",
+  "E72",
+  "E112",   # no-indented-block
+  "E113",   # unexpected-indentation
+  # "E124",
+  "E203",   # whitespace-before-punctuation
+  "E272",   # multiple-spaces-before-keyword
+  "E275",   # missing-whitespace-after-keyword
+  "E303",   # too-many-blank-lines
+  "E304",   # blank-line-after-decorator
+  "E501",   # line-too-long
+  # "E502",
+  "E702",   # multiple-statements-on-one-line-semicolon
+  "E703",   # useless-semicolon
+  "E731",   # lambda-assignment
+  "W191",   # tab-indentation
+  "W291",   # trailing-whitespace
+  "W293",   # blank-line-with-whitespace
+  "UP039",  # unnecessary-class-parentheses
+  "C416",   # unnecessary-comprehension
+  "RET506", # superfluous-else-raise
+  "RET507", # superfluous-else-continue
+  "A",      # builtin-variable-shadowing, builtin-argument-shadowing, builtin-attribute-shadowing
+  "FURB110",# if-exp-instead-of-or-operator
+  "RUF018", # assignment-in-assert
+]
+
+# detect unused imports in examples
+[tool.ruff.lint.per-file-ignores]
+"examples/**/*.py" = [
+  "W6",
+  "E71",
+  "E72",
+  "E112",
+  "E113",
+  "E203",
+  "E272",
+  "E275",
+  "E303",
+  "E304",
+  "E501",
+  "E702",
+  "E703",
+  "E731",
+  "W191",
+  "W291",
+  "W293",
+  "UP039",
+  "C416",
+  "RET506",
+  "RET507",
+  "A",
+  "FURB110",
+  "RUF018",
+  "F541",
+  "F841",
+]
+
+[tool.ruff.format]
+exclude = ["*"]
@@ -1,9 +0,0 @@
-[pytest]
-norecursedirs =
-    extra
-    .hypothesis
-    .git
-timeout = 300
-timeout_method = thread
-timeout_func_only = true
-testpaths = test
@@ -1,56 +0,0 @@
-indent-width = 2
-preview = true
-target-version = "py311"
-
-lint.select = [
-  "F",  # Pyflakes
-  "W6",
-  "E71",
-  "E72",
-  "E112",   # no-indented-block
-  "E113",   # unexpected-indentation
-  # "E124",
-  "E203",   # whitespace-before-punctuation
-  "E272",   # multiple-spaces-before-keyword
-  "E275",   # missing-whitespace-after-keyword
-  "E303",   # too-many-blank-lines
-  "E304",   # blank-line-after-decorator
-  "E501",   # line-too-long
-  # "E502",
-  "E702",   # multiple-statements-on-one-line-semicolon
-  "E703",   # useless-semicolon
-  "E731",   # lambda-assignment
-  "W191",   # tab-indentation
-  "W291",   # trailing-whitespace
-  "W293",   # blank-line-with-whitespace
-  "UP039",  # unnecessary-class-parentheses
-  "C416",   # unnecessary-comprehension
-  "RET506", # superfluous-else-raise
-  "RET507", # superfluous-else-continue
-  "A",      # builtin-variable-shadowing, builtin-argument-shadowing, builtin-attribute-shadowing
-  "FURB110",# if-exp-instead-of-or-operator
-  "RUF018", # assignment-in-assert
-]
-
-line-length = 150
-
-exclude = [
-  ".git/",
-  "docs/",
-  "extra/",
-  "tinygrad/runtime/autogen",
-  "test/external/mlperf_resnet",
-  "test/external/mlperf_unet3d",
-]
-
-# detect unused imports in examples
-[lint.per-file-ignores]
-"examples/**/*.py" = [
-  "W6", "E71", "E72", "E112", "E113", "E203", "E272", "E275",
-  "E303", "E304", "E501", "E702", "E703", "E731", "W191",
-  "W291", "W293", "UP039", "C416", "RET506", "RET507", "A",
-  "FURB110", "RUF018", "F541", "F841"
-]
-
-[format]
-exclude = ["*"]
@@ -1,21 +0,0 @@
-[mutmut]
-paths_to_mutate=tinygrad
-do_not_mutate=
-  tinygrad/apps/*
-  tinygrad/codegen/*
-  tinygrad/engine/*
-  tinygrad/nn/*
-  tinygrad/renderer/*
-  tinygrad/runtime/*
-  tinygrad/schedule/*
-  tinygrad/uop/*
-  tinygrad/viz/*
-  tinygrad/device.py
-  tinygrad/dtype.py
-  tinygrad/gradient.py
-  tinygrad/helpers.py
-  tinygrad/tensor.py
-tests_dir=
-  test/test_tiny.py
-  test/test_ops.py
-debug=true
@@ -1,111 +0,0 @@
-#!/usr/bin/env python3
-
-from pathlib import Path
-from setuptools import setup
-
-directory = Path(__file__).resolve().parent
-with open(directory / 'README.md', encoding='utf-8') as f:
-  long_description = f.read()
-
-testing_minimal = [
-  "numpy",
-  "torch==2.9.0",
-  "pytest",
-  "pytest-xdist",
-  "pytest-timeout",
-  "pytest-split",
-  "hypothesis",
-  "z3-solver",
-]
-
-setup(name='tinygrad',
-      version='0.11.0',
-      description='You like pytorch? You like micrograd? You love tinygrad! <3',
-      author='George Hotz',
-      license='MIT',
-      long_description=long_description,
-      long_description_content_type='text/markdown',
-      packages = [
-        'tinygrad',
-        'tinygrad.apps',
-        'tinygrad.codegen',
-        'tinygrad.codegen.opt',
-        'tinygrad.codegen.late',
-        'tinygrad.engine',
-        'tinygrad.mixin',
-        'tinygrad.nn',
-        'tinygrad.renderer',
-        'tinygrad.runtime',
-        'tinygrad.runtime.autogen',
-        'tinygrad.runtime.autogen.am',
-        'tinygrad.runtime.autogen.nv',
-        'tinygrad.runtime.graph',
-        'tinygrad.runtime.support',
-        'tinygrad.runtime.support.am',
-        'tinygrad.runtime.support.nv',
-        'tinygrad.schedule',
-        'tinygrad.uop',
-        'tinygrad.viz',
-      ],
-      package_data = {'tinygrad': ['py.typed'], 'tinygrad.viz': ['index.html', 'assets/**/*', 'js/*']},
-      classifiers=[
-        "Programming Language :: Python :: 3",
-        "License :: OSI Approved :: MIT License"
-      ],
-      install_requires=[],
-      python_requires='>=3.11',
-      extras_require={
-        'arm': ["unicorn"],
-        'triton': ["triton-nightly>=2.1.0.dev20231014192330"],
-        'linting': [
-            "pylint",
-            "mypy==1.18.1",
-            "typing-extensions",
-            "pre-commit",
-            "ruff",
-            "numpy",
-            "typeguard",
-        ],
-        #'mlperf': ["mlperf-logging @ git+https://github.com/mlperf/[email protected]"],
-        'testing_minimal': testing_minimal,
-        'testing_unit': testing_minimal + [
-            "tqdm",
-            "safetensors",
-            "tabulate",  # for sz.py
-        ],
-        'testing': testing_minimal + [
-            "pillow",
-            "onnx==1.18.0",
-            "onnx2torch",
-            "onnxruntime",
-            "opencv-python",
-            "tabulate",
-            "tqdm",
-            "safetensors",
-            "transformers",
-            "sentencepiece",
-            "tiktoken",
-            "blobfile",
-            "librosa",
-            "numba>=0.55",  # librosa needs numba but uv ignores python upper bounds and some numba versions require <python3.10
-            "networkx",
-            "nibabel",
-            "bottle",
-            "ggml-python",
-            "capstone",
-            "pycocotools",
-            "boto3",
-            "pandas",
-            "influxdb3-python"
-        ],
-        'docs': [
-            "mkdocs",
-            "mkdocs-material",
-            "mkdocstrings[python]",
-            "markdown-callouts",
-            "markdown-exec[ansi]",
-            "black",
-            "numpy",
-        ],
-      },
-      include_package_data=True)
@@ -1,163 +0,0 @@
-# ruff: noqa: E501 E712
-from tinygrad import dtypes, Device
-from tinygrad.uop.ops import UOp, AxisType, Ops, KernelInfo
-from tinygrad.codegen import full_rewrite
-from tinygrad.renderer import ProgramSpec
-from tinygrad.engine.realize import CompiledRunner
-from tinygrad.helpers import dedup
-from tinygrad.device import Buffer
-from tinygrad.dtype import ImageDType, Invalid
-
-c0 = UOp(Ops.DEFINE_GLOBAL, dtypes.float.ptr(1576), (), 0)
-c2 = UOp.range(1576, 20, AxisType.LOOP)
-c5 = c2<55
-c6 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((1, 16, 4)), (), 1)
-c8 = UOp.range(16, 0, AxisType.REDUCE)
-c11 = UOp.range(4, 1, AxisType.REDUCE)
-c14 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((14, 64, 4)), (), 2)
-c25 = c5.where((c2%4*4+c11+c8*16+c2//4*256), UOp.const(dtypes.index, Invalid))
-c27 = c6.index((c8*4+c11))*c14.index(c25)
-c29 = UOp(Ops.DEFINE_GLOBAL, dtypes.float.ptr(55), (), 3)
-c30 = c5.where(c2, UOp.const(dtypes.index, Invalid))
-c34 = c5.where((c27.reduce(c8, c11, arg=Ops.ADD)+c29.index(c30)), UOp.const(dtypes.float, 0.0))
-c38 = c2<87
-c39 = (c5!=True)&c38
-c40 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((1, 8, 4)), (), 4)
-c42 = UOp.range(8, 2, AxisType.REDUCE)
-c44 = UOp.range(4, 3, AxisType.REDUCE)
-c47 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((8, 32, 4)), (), 5)
-c49 = c2+1
-c51 = c49%4*4
-c57 = c49//4*128
-c61 = c39.where((c51+c44+c42*16+c57+-1792), UOp.const(dtypes.index, Invalid))
-c63 = c40.index((c42*4+c44))*c47.index(c61)
-c65 = UOp(Ops.DEFINE_GLOBAL, dtypes.float.ptr(32), (), 6)
-c68 = c39.where((c2+-55), UOp.const(dtypes.index, Invalid))
-c71 = c39.where((c63.reduce(c42, c44, arg=Ops.ADD)+c65.index(c68)), UOp.const(dtypes.float, 0.0))
-c75 = c2<99
-c76 = (c38!=True)&c75
-c77 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((1, 8, 4)), (), 7)
-c78 = UOp.range(8, 4, AxisType.REDUCE)
-c80 = UOp.range(4, 5, AxisType.REDUCE)
-c83 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((3, 32, 4)), (), 8)
-c90 = c76.where((c51+c80+c78*16+c57+-2816), UOp.const(dtypes.index, Invalid))
-c92 = c77.index((c78*4+c80))*c83.index(c90)
-c94 = UOp(Ops.DEFINE_GLOBAL, dtypes.float.ptr(12), (), 9)
-c97 = c76.where((c2+-87), UOp.const(dtypes.index, Invalid))
-c100 = c76.where((c92.reduce(c78, c80, arg=Ops.ADD)+c94.index(c97)), UOp.const(dtypes.float, 0.0))
-c104 = c2<105
-c105 = (c75!=True)&c104
-c106 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((1, 8, 4)), (), 10)
-c107 = UOp.range(8, 6, AxisType.REDUCE)
-c109 = UOp.range(4, 7, AxisType.REDUCE)
-c112 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((2, 32, 4)), (), 11)
-c119 = c105.where((c51+c109+c107*16+c57+-3200), UOp.const(dtypes.index, Invalid))
-c121 = c106.index((c107*4+c109))*c112.index(c119)
-c123 = UOp(Ops.DEFINE_GLOBAL, dtypes.float.ptr(6), (), 12)
-c126 = c105.where((c2+-99), UOp.const(dtypes.index, Invalid))
-c129 = c105.where((c121.reduce(c107, c109, arg=Ops.ADD)+c123.index(c126)), UOp.const(dtypes.float, 0.0))
-c133 = c2<117
-c134 = (c104!=True)&c133
-c135 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((1, 8, 4)), (), 13)
-c136 = UOp.range(8, 8, AxisType.REDUCE)
-c138 = UOp.range(4, 9, AxisType.REDUCE)
-c141 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((3, 32, 4)), (), 14)
-c143 = c2+3
-c145 = c143%4*4
-c149 = c143//4
-c150 = c149*128
-c154 = c134.where((c145+c138+c136*16+c150+-3456), UOp.const(dtypes.index, Invalid))
-c156 = c135.index((c136*4+c138))*c141.index(c154)
-c158 = UOp(Ops.DEFINE_GLOBAL, dtypes.float.ptr(12), (), 15)
-c161 = c134.where((c2+-105), UOp.const(dtypes.index, Invalid))
-c164 = c134.where((c156.reduce(c136, c138, arg=Ops.ADD)+c158.index(c161)), UOp.const(dtypes.float, 0.0))
-c168 = c2<645
-c169 = (c133!=True)&c168
-c170 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((1, 16, 4)), (), 16)
-c171 = UOp.range(16, 10, AxisType.REDUCE)
-c173 = UOp.range(4, 11, AxisType.REDUCE)
-c176 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((132, 64, 4)), (), 17)
-c180 = c149*256
-c184 = c169.where((c145+c173+c171*16+c180+-7680), UOp.const(dtypes.index, Invalid))
-c186 = c170.index((c171*4+c173))*c176.index(c184)
-c188 = UOp(Ops.DEFINE_GLOBAL, dtypes.float.ptr(528), (), 18)
-c191 = c169.where((c2+-117), UOp.const(dtypes.index, Invalid))
-c194 = c169.where((c186.reduce(c171, c173, arg=Ops.ADD)+c188.index(c191)), UOp.const(dtypes.float, 0.0))
-c198 = c2<653
-c199 = (c168!=True)&c198
-c200 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((1, 4, 4)), (), 19)
-c201 = UOp.range(4, 12, AxisType.REDUCE)
-c203 = UOp.range(4, 13, AxisType.REDUCE)
-c206 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((2, 16, 4)), (), 20)
-c215 = c199.where((c145+c203+c201*16+c149*64+-10368), UOp.const(dtypes.index, Invalid))
-c217 = c200.index((c201*4+c203))*c206.index(c215)
-c219 = UOp(Ops.DEFINE_GLOBAL, dtypes.float.ptr(8), (), 21)
-c222 = c199.where((c2+-645), UOp.const(dtypes.index, Invalid))
-c225 = c199.where((c217.reduce(c201, c203, arg=Ops.ADD)+c219.index(c222)), UOp.const(dtypes.float, 0.0))
-c229 = c2<917
-c230 = (c198!=True)&c229
-c231 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((1, 8, 4)), (), 22)
-c232 = UOp.range(8, 14, AxisType.REDUCE)
-c234 = UOp.range(4, 15, AxisType.REDUCE)
-c237 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((66, 32, 4)), (), 23)
-c244 = c230.where((c145+c234+c232*16+c150+-20992), UOp.const(dtypes.index, Invalid))
-c246 = c231.index((c232*4+c234))*c237.index(c244)
-c248 = UOp(Ops.DEFINE_GLOBAL, dtypes.float.ptr(264), (), 24)
-c251 = c230.where((c2+-653), UOp.const(dtypes.index, Invalid))
-c254 = c230.where((c246.reduce(c232, c234, arg=Ops.ADD)+c248.index(c251)), UOp.const(dtypes.float, 0.0))
-c258 = c2<1061
-c259 = (c229!=True)&c258
-c260 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((1, 16, 4)), (), 25)
-c261 = UOp.range(16, 16, AxisType.REDUCE)
-c263 = UOp.range(4, 17, AxisType.REDUCE)
-c266 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((36, 64, 4)), (), 26)
-c273 = c259.where((c145+c263+c261*16+c180+-58880), UOp.const(dtypes.index, Invalid))
-c275 = c260.index((c261*4+c263))*c266.index(c273)
-c277 = UOp(Ops.DEFINE_GLOBAL, dtypes.float.ptr(144), (), 27)
-c280 = c259.where((c2+-917), UOp.const(dtypes.index, Invalid))
-c283 = UOp(Ops.DEFINE_GLOBAL, dtypes.float.ptr(144), (), 28)
-c286 = c259.where(((c275.reduce(c261, c263, arg=Ops.ADD)+c277.index(c280))*c283.index(c280)), UOp.const(dtypes.float, 0.0))
-c290 = c2<1064
-c291 = (c258!=True)&c290
-c292 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((1, 4, 4)), (), 29)
-c293 = UOp.range(4, 18, AxisType.REDUCE)
-c295 = UOp.range(4, 19, AxisType.REDUCE)
-c298 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((1, 16, 4)), (), 30)
-c305 = c291.where((c2*4+c295+c293*16+-4244), UOp.const(dtypes.index, Invalid))
-c307 = c292.index((c293*4+c295))*c298.index(c305)
-c309 = UOp(Ops.DEFINE_GLOBAL, dtypes.float.ptr(3), (), 31)
-c312 = c291.where((c2+-1061), UOp.const(dtypes.index, Invalid))
-c315 = c291.where((c307.reduce(c293, c295, arg=Ops.ADD)+c309.index(c312)), UOp.const(dtypes.float, 0.0))
-c317 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((1, 128, 4)), (), 32)
-c321 = (c290!=True).where((c2+-1064), UOp.const(dtypes.index, Invalid))
-c323 = UOp(Ops.DEFINE_GLOBAL, dtypes.float.ptr(1), (), 33)
-c328 = c290.where(UOp.const(dtypes.float, 0.0), (c317.index(c321)*c323.index(UOp.const(dtypes.index, 0)).reciprocal()))
-c329 = c34+c71+c100+c129+c164+c194+c225+c254+c286+c315+c328
-c331 = c0.index(c2, ptr=True).store(c329).end(c2)
-ast = c331.sink(arg=KernelInfo(name="cat", opts_to_apply=None))
-
-compiler = Device.default.compiler
-renderer = Device.default.renderer
-allocator = Device.default.allocator
-
-uops = full_rewrite(ast, renderer)
-src = renderer.render(uops)
-
-# NOLOCALS=1 IMAGE=2 DEV=CL
-lib = compiler.compile(src)
-
-ps = ProgramSpec("cat", src, Device.DEFAULT, ast, uops)
-# print(ps.src)
-# print(ps.applied_opts)
-# NOTE: this is faster with no GROUP and with NOLOCALS
-# (Opt(op=OptOps.UPCAST, axis=1, arg=4), Opt(op=OptOps.UNROLL, axis=19, arg=4), Opt(op=OptOps.UNROLL, axis=17, arg=4), Opt(op=OptOps.UNROLL, axis=15, arg=4), Opt(op=OptOps.UNROLL, axis=13, arg=4), Opt(op=OptOps.UNROLL, axis=11, arg=4), Opt(op=OptOps.UNROLL, axis=9, arg=4), Opt(op=OptOps.UNROLL, axis=7, arg=4), Opt(op=OptOps.UNROLL, axis=5, arg=4), Opt(op=OptOps.UNROLL, axis=3, arg=4), Opt(op=OptOps.UNROLL, axis=1, arg=4), Opt(op=OptOps.NOLOCALS, axis=None, arg=None))
-cr = CompiledRunner(ps, precompiled=lib)
-
-gs = sorted(dedup([u for u in ast.toposort() if u.op is Ops.DEFINE_GLOBAL]), key=lambda u: u.arg)
-print(len(gs))
-print([g.dtype for g in gs])
-
-bufs = [Buffer(ps.device, g.size, g.dtype if isinstance(g.dtype, ImageDType) else g.dtype._base).ensure_allocated() for g in gs]
-
-t = cr(bufs, wait=True)
-print(f"{t*1e6:.2f} us")
@@ -1,8 +1,8 @@
-# ruff: noqa: E501 E712
+# ruff: noqa: E501 E712 F401
 from tinygrad import dtypes, Device
 from tinygrad.uop.ops import UOp, AxisType, Ops, KernelInfo
 from tinygrad.codegen import full_rewrite
-# from tinygrad.codegen.opt import Opt, OptOps
+from tinygrad.codegen.opt import Opt, OptOps # pylint: disable=unused-import
 from tinygrad.renderer import ProgramSpec
 from tinygrad.engine.realize import CompiledRunner
 from tinygrad.helpers import dedup, getenv
@@ -33,6 +33,8 @@ def vision_conv_143():
  c67 = c0.index((c2*128+c5+c8*4096), ptr=True).store(c65).end(c8, c2, c5)

  opts = None
+  # JITBEAM=2
+  # (Opt(op=OptOps.UPCAST, axis=2, arg=4), Opt(op=OptOps.NOLOCALS, axis=None, arg=None), Opt(op=OptOps.UPCAST, axis=2, arg=2), Opt(op=OptOps.UPCAST, axis=1, arg=4), Opt(op=OptOps.SWAP, axis=1, arg=2))
  return c67.sink(arg=KernelInfo(name="conv", opts_to_apply=opts))

 def vision_conv_153():
@@ -57,9 +59,32 @@ def vision_conv_153():
  c67 = c0.index((c2*256+c5+c8*4096), ptr=True).store(c65).end(c8, c2, c5)

  opts = None
+  # JITBEAM=2
+  # (Opt(op=OptOps.UPCAST, axis=2, arg=4), Opt(op=OptOps.NOLOCALS, axis=None, arg=None), Opt(op=OptOps.UPCAST, axis=2, arg=2), Opt(op=OptOps.SWAP, axis=1, arg=2))
  return c67.sink(arg=KernelInfo(name="conv", opts_to_apply=opts))

-ast = vision_conv_143() if getenv("NUM", 143) == 143 else vision_conv_153()
+def dm_conv_172():
+  c0 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((1, 240, 4)), (), 0)
+  c2 = UOp.range(960, 4, AxisType.LOOP)
+  c5 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((8, 384, 4)), (), 1)
+  c7 = UOp.range(32, 0, AxisType.REDUCE)
+  c10 = UOp.range(4, 1, AxisType.REDUCE)
+  c13 = UOp.range(12, 3, AxisType.REDUCE)
+  c18 = UOp.range(8, 2, AxisType.REDUCE)
+  c23 = UOp(Ops.DEFINE_GLOBAL, dtypes.imageh((240, 128, 4)), (), 2)
+  c35 = c5.index((c7*4+c10+c13*128+c18*1536))*c23.index((c10*4+c2%4+c7*16+c2//4*512))
+  c37 = UOp(Ops.DEFINE_GLOBAL, dtypes.float.ptr(960), (), 3)
+  c39 = c35.reduce(c7, c10, arg=Ops.ADD)+c37.index(c2)
+  c50 = (1.0+((c39+0.044708251953125*(c39*(c39*c39)))*-2.3021129851685216).exp2()).reciprocal()*c39
+  c53 = c50.reduce(c18, c13, arg=Ops.ADD)*0.010416666666666666
+  c55 = c0.index(c2, ptr=True).store(c53).end(c2)
+
+  opts = None
+  # JITBEAM=2
+  # (Opt(op=OptOps.UPCAST, axis=0, arg=4), Opt(op=OptOps.GROUPTOP, axis=1, arg=32), Opt(op=OptOps.UNROLL, axis=1, arg=4), Opt(op=OptOps.LOCAL, axis=0, arg=8), Opt(op=OptOps.UNROLL, axis=0, arg=4), Opt(op=OptOps.GROUP, axis=1, arg=0))
+  return c55.sink(arg=KernelInfo(name="conv", opts_to_apply=opts))
+
+ast = {143: vision_conv_143, 153: vision_conv_153, 172: dm_conv_172}[getenv("NUM", 143)]()

 compiler = Device.default.compiler
 renderer = Device.default.renderer
@@ -112,7 +112,6 @@ backend_test.exclude('test_dequantizelinear_e5m2_cpu')
 backend_test.exclude('test_dequantizelinear_float4e2m1_cpu')

 # we don't support indexes
-backend_test.exclude('test_nonzero_*')

 # no support for int pow
 backend_test.exclude('test_pow_types_int32_int32_cpu')
@@ -78,6 +78,7 @@ class TestLinearizer(unittest.TestCase):
    ranges = [i for i,u in enumerate(uops) if u.op is Ops.RANGE]
    assert len(ranges) == 1 # NOTE: it collapses now

+  @unittest.expectedFailure # TODO: investigate
  def test_two_nested_range_alt_indexing(self):
    a = Tensor([2, 2]).realize()
    out = a.reshape(2, 1).pad(((1, 1), (1, 1)), value=2).sum()
@@ -1,18 +0,0 @@
-from tinygrad import Tensor, UOp
-from tinygrad.uop.ops import Ops, AxisType
-import unittest
-# this test is only focused on transformers and using range for the layers
-
-class TestOuterworldTransformer(unittest.TestCase):
-  def test_three_mats(self):
-    w = Tensor.empty(3, 1024, 1024)
-    inp = Tensor.empty(1, 1024)
-    i = UOp.range(3, -1, AxisType.OUTER)
-    inp_after = Tensor(inp.uop.after(i))
-    inp_gemm = inp_after@w[i]
-    inp = inp.uop.after(inp.uop.store(inp_gemm.uop).end(i)).contiguous()
-    inp = Tensor(inp)
-    inp.realize()
-
-if __name__ == "__main__":
-  unittest.main()
@@ -3,7 +3,7 @@ import torch
 import unittest, copy, mmap, random, math, array
 from tinygrad import Tensor, Device, dtypes
 from tinygrad.tensor import _METADATA
-from tinygrad.helpers import getenv, temp, mv_address
+from tinygrad.helpers import Context, getenv, temp, mv_address
 from extra.gradcheck import numerical_jacobian, jacobian, gradcheck
 from hypothesis import given, settings, strategies as strat
 from tinygrad.device import is_dtype_supported
@@ -846,6 +846,16 @@ class TestTensorMetadata(unittest.TestCase):
    #self.assertEqual(len(bw), 1)
    #self.assertEqual(bw[0].name, "sigmoid")

+  def test_tracemeta_0(self):
+    with Context(TRACEMETA=0):
+      x = Tensor.rand(3, requires_grad=True)
+      y = Tensor.rand(3, requires_grad=True)
+      out = (x.relu() * y.sigmoid()).sum()
+      self.assertIsNone(out.uop.metadata)
+      self.assertIsNone(out.uop.src[0].metadata)
+      si = out.schedule()[-1]
+      self.assertEqual(si.metadata, ())
+
 class TestIdxUpcast(unittest.TestCase):
  def _find_op(self, ast: UOp, op: Ops):
    if ast.op is op: return ast
@@ -1,31 +1,35 @@
-import unittest
+import unittest, math

 from tinygrad import Tensor, Device, dtypes, Context
 from tinygrad.engine.realize import ExecItem, get_runner
+from tinygrad.helpers import CI
+from tinygrad.renderer.ptx import PTXRenderer
+import numpy as np

 from extra.thunder.tiny.tk import WARP_THREADS
 from extra.thunder.tiny.tk.kernel import Kernel
-from extra.thunder.tiny.tk.tiles import gl, st, rt, rv

+@unittest.skipUnless(Device.DEFAULT in ["CUDA", "NV"], "only cuda")
+@unittest.skipIf(isinstance(Device[Device.DEFAULT].renderer, PTXRenderer), "no ptx")
 class TestTK(unittest.TestCase):
-  @unittest.skip("store from float rt is wrong")
+  @unittest.skipIf(CI, "no wmma in ci")
  def test_simple_matmul(self):
    N = 32
    BLOCK_SIZE = 16
    with Kernel((N // BLOCK_SIZE, N // BLOCK_SIZE, 1), WARP_THREADS) as ker:
      warp = ker.warp

-      c = gl((1, 1, N, N), dtypes.float32)
-      a = gl((1, 1, N, N), dtypes.bfloat16)
-      b = gl((1, 1, N, N), dtypes.bfloat16)
+      c = ker.gl((1, 1, N, N), dtypes.float32)
+      a = ker.gl((1, 1, N, N), dtypes.bfloat16)
+      b = ker.gl((1, 1, N, N), dtypes.bfloat16)

-      a_smem = st((BLOCK_SIZE, BLOCK_SIZE), dtypes.bfloat16)
-      b_smem = st((BLOCK_SIZE, BLOCK_SIZE), dtypes.bfloat16)
-      c_smem = st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+      a_smem = ker.st((BLOCK_SIZE, BLOCK_SIZE), dtypes.bfloat16)
+      b_smem = ker.st((BLOCK_SIZE, BLOCK_SIZE), dtypes.bfloat16)
+      c_smem = ker.st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)

-      a_reg = rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.bfloat16)
-      b_reg = rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.bfloat16)
-      c_reg = rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+      a_reg = ker.rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.bfloat16)
+      b_reg = ker.rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.bfloat16)
+      c_reg = ker.rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)

      col, row = ker.blockIdx_x, ker.blockIdx_y

@@ -57,26 +61,26 @@ class TestTK(unittest.TestCase):

    ref = a.matmul(b, dtype=dtypes.float32).float()

-    assert ref.allclose(c)
+    np.testing.assert_allclose(c.numpy(), ref.numpy())

-  @unittest.skip("store from float rt is wrong")
+  @unittest.skipIf(CI, "no wmma in ci")
  def test_simple_matmul_transposed(self):
    N = 32
    BLOCK_SIZE = 16
    with Kernel((N // BLOCK_SIZE, N // BLOCK_SIZE, 1), WARP_THREADS) as ker:
      warp = ker.warp

-      c = gl((1, 1, N, N), dtypes.float32)
-      a = gl((1, 1, N, N), dtypes.bfloat16)
-      b = gl((1, 1, N, N), dtypes.bfloat16)
+      c = ker.gl((1, 1, N, N), dtypes.float32)
+      a = ker.gl((1, 1, N, N), dtypes.bfloat16)
+      b = ker.gl((1, 1, N, N), dtypes.bfloat16)

-      a_smem = st((BLOCK_SIZE, BLOCK_SIZE), dtypes.bfloat16)
-      b_smem = st((BLOCK_SIZE, BLOCK_SIZE), dtypes.bfloat16)
-      c_smem = st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+      a_smem = ker.st((BLOCK_SIZE, BLOCK_SIZE), dtypes.bfloat16)
+      b_smem = ker.st((BLOCK_SIZE, BLOCK_SIZE), dtypes.bfloat16)
+      c_smem = ker.st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)

-      a_reg = rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.bfloat16)
-      b_reg = rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.bfloat16)
-      c_reg = rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+      a_reg = ker.rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.bfloat16)
+      b_reg = ker.rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.bfloat16)
+      c_reg = ker.rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)

      col, row = ker.blockIdx_x, ker.blockIdx_y

@@ -108,7 +112,7 @@ class TestTK(unittest.TestCase):

    ref = a.matmul(b.transpose(2, 3), dtype=dtypes.float32).float()

-    assert ref.allclose(c)
+    np.testing.assert_allclose(c.numpy(), ref.numpy())

  def test_load_store(self):
    N = 32
@@ -116,14 +120,14 @@ class TestTK(unittest.TestCase):
    with Kernel((N // BLOCK_SIZE, N // BLOCK_SIZE, 1), WARP_THREADS) as ker:
      warp = ker.warp

-      b = gl((1, 1, N, N), dtypes.float32)
-      a = gl((1, 1, N, N), dtypes.float32)
+      b = ker.gl((1, 1, N, N), dtypes.float32)
+      a = ker.gl((1, 1, N, N), dtypes.float32)

-      a_smem = st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
-      b_smem = st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+      a_smem = ker.st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+      b_smem = ker.st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)

-      a_reg = rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
-      b_reg = rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+      a_reg = ker.rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+      b_reg = ker.rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)

      col, row = ker.blockIdx_x, ker.blockIdx_y

@@ -146,7 +150,7 @@ class TestTK(unittest.TestCase):

    ref = a.float()

-    assert ref.allclose(b)
+    np.testing.assert_allclose(b.numpy(), ref.numpy())

  def test_max(self):
    N = 16
@@ -154,28 +158,27 @@ class TestTK(unittest.TestCase):
    with Kernel((1, 1, 1), WARP_THREADS) as ker:
      warp = ker.warp

-      b = gl((1, 1, N, N), dtypes.float32)
-      a = gl((1, 1, N, N), dtypes.float32)
+      b = ker.gl((1, 1, N, N), dtypes.float32)
+      a = ker.gl((1, 1, N, N), dtypes.float32)

-      a_smem = st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
-      b_smem = st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+      a_smem = ker.st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+      b_smem = ker.st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)

-      a_reg = rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
-      b_reg = rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+      a_reg = ker.rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+      b_reg = ker.rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)

-      max_reg = rv(BLOCK_SIZE, dtypes.float32, "ortho")
-
-      max_reg = warp.neg_inf(max_reg)
+      max_reg = ker.rv(BLOCK_SIZE, dtypes.float32, "ortho")

      for tile_row in ker.range(N // BLOCK_SIZE):
+        max_reg = warp.neg_inf(max_reg.after(tile_row))
+
        for tile_col in ker.range(N // BLOCK_SIZE):
          a_smem = warp.load(a_smem, a, (), (0, 0, tile_row, tile_col), axis=2)
          a_reg = warp.load(a_reg, a_smem)
          max_reg = warp.row_reduce(max_reg, a_reg, lambda a, b: a.maximum(b))
-        sum_reg = ker.endrange()
+        max_reg = ker.endrange()

-        b_reg = warp.zero(b_reg).after(tile_row)
-        b_reg = warp.map(b_reg, lambda _, idx: sum_reg[idx[0], 0, (idx[2]%4)//2])
+        b_reg = warp.map(b_reg, lambda _, idx: max_reg[idx[0], 0, (idx[2]%4)//2])
        b_smem = warp.store(b_smem, b_reg)

        for tile_col in ker.range(N // BLOCK_SIZE):
@@ -194,7 +197,7 @@ class TestTK(unittest.TestCase):

    ref = a.float().max(axis=3, keepdim=True).expand(a.shape)

-    assert ref.allclose(b)
+    np.testing.assert_allclose(b.numpy(), ref.numpy())

  def test_max_nonsquare(self):
    N, M = 16, 64
@@ -202,28 +205,27 @@ class TestTK(unittest.TestCase):
    with Kernel((1, 1, 1), WARP_THREADS) as ker:
      warp = ker.warp

-      b = gl((1, 1, N, M), dtypes.float32)
-      a = gl((1, 1, N, M), dtypes.float32)
+      b = ker.gl((1, 1, N, M), dtypes.float32)
+      a = ker.gl((1, 1, N, M), dtypes.float32)

-      a_smem = st((BLOCK_N, BLOCK_M), dtypes.float32)
-      b_smem = st((BLOCK_N, BLOCK_M), dtypes.float32)
+      a_smem = ker.st((BLOCK_N, BLOCK_M), dtypes.float32)
+      b_smem = ker.st((BLOCK_N, BLOCK_M), dtypes.float32)

-      a_reg = rt((BLOCK_N, BLOCK_M), dtypes.float32)
-      b_reg = rt((BLOCK_N, BLOCK_M), dtypes.float32)
+      a_reg = ker.rt((BLOCK_N, BLOCK_M), dtypes.float32)
+      b_reg = ker.rt((BLOCK_N, BLOCK_M), dtypes.float32)

-      max_reg = rv(BLOCK_N, dtypes.float32, "ortho")
-
-      max_reg = warp.zero(max_reg)
+      max_reg = ker.rv(BLOCK_N, dtypes.float32, "ortho")

      for tile_row in ker.range(N // BLOCK_N):
+        max_reg = warp.neg_inf(max_reg.after(tile_row))
+
        for tile_col in ker.range(M // BLOCK_M):
          a_smem = warp.load(a_smem, a, (), (0, 0, tile_row, tile_col), axis=2)
          a_reg = warp.load(a_reg, a_smem)
-          sum_reg = warp.row_reduce(max_reg, a_reg, lambda a, b: a.maximum(b))
-        sum_reg = ker.endrange()
+          max_reg = warp.row_reduce(max_reg, a_reg, lambda a, b: a.maximum(b))
+        max_reg = ker.endrange()

-        b_reg = warp.zero(b_reg).after(tile_row)
-        b_reg = warp.map(b_reg, lambda _, idx: sum_reg[idx[0], 0, (idx[2]%4)//2])
+        b_reg = warp.map(b_reg, lambda _, idx: max_reg[idx[0], 0, (idx[2]%4)//2])
        b_smem = warp.store(b_smem, b_reg)

        for tile_col in ker.range(M // BLOCK_M):
@@ -242,27 +244,27 @@ class TestTK(unittest.TestCase):

    ref = a.float().max(axis=3, keepdim=True).expand(a.shape)

-    assert ref.allclose(b)
+    np.testing.assert_allclose(b.numpy(), ref.numpy())

  def test_sum(self):
-    N = 16
+    N = 32
    BLOCK_SIZE = 16
    with Kernel((1, 1, 1), WARP_THREADS) as ker:
      warp = ker.warp

-      b = gl((1, 1, N, N), dtypes.float32)
-      a = gl((1, 1, N, N), dtypes.float32)
+      b = ker.gl((1, 1, N, N), dtypes.float32)
+      a = ker.gl((1, 1, N, N), dtypes.float32)

-      a_smem = st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
-      b_smem = st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+      a_smem = ker.st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+      b_smem = ker.st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)

-      a_reg = rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
-      b_reg = rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+      a_reg = ker.rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+      b_reg = ker.rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)

-      sum_reg = rv(BLOCK_SIZE, dtypes.float32, "ortho")
+      sum_reg = ker.rv(BLOCK_SIZE, dtypes.float32, "ortho")

      for tile_row in ker.range(N // BLOCK_SIZE):
-        sum_reg = warp.zero(sum_reg).after(tile_row)
+        sum_reg = warp.zero(sum_reg.after(tile_row))

        for tile_col in ker.range(N // BLOCK_SIZE):
          a_smem = warp.load(a_smem, a, (), (0, 0, tile_row, tile_col), axis=2)
@@ -270,7 +272,6 @@ class TestTK(unittest.TestCase):
          sum_reg = warp.row_reduce(sum_reg, a_reg, lambda a, b: a + b)
        sum_reg = ker.endrange()

-        b_reg = warp.zero(b_reg).after(tile_row)
        b_reg = warp.map(b_reg, lambda _, idx: sum_reg[idx[0], 0, (idx[2]%4)//2])
        b_smem = warp.store(b_smem, b_reg)

@@ -281,7 +282,6 @@ class TestTK(unittest.TestCase):

    with Context(DEBUG=0):
      a = Tensor.rand(1, 1, N, N, dtype="float32").contiguous()
-      a = Tensor.arange(1 * 1 * N * N).reshape(1, 1, N, N).cast(dtypes.float32).contiguous()
      b = Tensor.empty(1, 1, N, N, dtype="float32")
      Tensor.realize(a, b)

@@ -291,7 +291,7 @@ class TestTK(unittest.TestCase):

    ref = a.float().sum(axis=3, keepdim=True).expand(a.shape)

-    assert ref.allclose(b)
+    np.testing.assert_allclose(b.numpy(), ref.numpy(), atol=1e-5, rtol=1e-5)

  def test_sum_nonsquare(self):
    N, M = 16, 64
@@ -299,27 +299,26 @@ class TestTK(unittest.TestCase):
    with Kernel((1, 1, 1), WARP_THREADS) as ker:
      warp = ker.warp

-      b = gl((1, 1, N, M), dtypes.float32)
-      a = gl((1, 1, N, M), dtypes.float32)
+      b = ker.gl((1, 1, N, M), dtypes.float32)
+      a = ker.gl((1, 1, N, M), dtypes.float32)

-      a_smem = st((BLOCK_N, BLOCK_M), dtypes.float32)
-      b_smem = st((BLOCK_N, BLOCK_M), dtypes.float32)
+      a_smem = ker.st((BLOCK_N, BLOCK_M), dtypes.float32)
+      b_smem = ker.st((BLOCK_N, BLOCK_M), dtypes.float32)

-      a_reg = rt((BLOCK_N, BLOCK_M), dtypes.float32)
-      b_reg = rt((BLOCK_N, BLOCK_M), dtypes.float32)
+      a_reg = ker.rt((BLOCK_N, BLOCK_M), dtypes.float32)
+      b_reg = ker.rt((BLOCK_N, BLOCK_M), dtypes.float32)

-      sum_reg = rv(BLOCK_N, dtypes.float32, "ortho")
-
-      sum_reg = warp.zero(sum_reg)
+      sum_reg = ker.rv(BLOCK_N, dtypes.float32, "ortho")

      for tile_row in ker.range(N // BLOCK_N):
+        sum_reg = warp.zero(sum_reg.after(tile_row))
+
        for tile_col in ker.range(M // BLOCK_M):
          a_smem = warp.load(a_smem, a, (), (0, 0, tile_row, tile_col), axis=2)
          a_reg = warp.load(a_reg, a_smem)
          sum_reg = warp.row_reduce(sum_reg, a_reg, lambda a, b: a + b)
        sum_reg = ker.endrange()

-        b_reg = warp.zero(b_reg).after(tile_row)
        b_reg = warp.map(b_reg, lambda _, idx: sum_reg[idx[0], 0, (idx[2]%4)//2])
        b_smem = warp.store(b_smem, b_reg)

@@ -339,7 +338,68 @@ class TestTK(unittest.TestCase):

    ref = a.float().sum(axis=3, keepdim=True).expand(a.shape)

-    assert ref.allclose(b)
+    np.testing.assert_allclose(b.numpy(), ref.numpy(), atol=1e-5, rtol=1e-5)
+
+  @unittest.skip("fake range not ended")
+  def test_softmax(self):
+    N = 32
+    BLOCK_SIZE = 16
+    with Kernel((1, 1, 1), WARP_THREADS) as ker:
+      warp = ker.warp
+
+      b = ker.gl((1, 1, BLOCK_SIZE, N), dtypes.float32)
+      a = ker.gl((1, 1, BLOCK_SIZE, N), dtypes.float32)
+
+      a_smem = ker.st((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+
+      a_reg = ker.rt((BLOCK_SIZE, BLOCK_SIZE), dtypes.float32)
+
+      max_vec_last = ker.rv(BLOCK_SIZE, dtypes.float32, "ortho")
+      max_vec = ker.rv(BLOCK_SIZE, dtypes.float32, "ortho")
+      norm_vec = ker.rv(BLOCK_SIZE, dtypes.float32, "ortho")
+
+      max_vec = warp.neg_inf(max_vec)
+      norm_vec = warp.zero(norm_vec)
+
+      for tile_col in ker.range(N // BLOCK_SIZE):
+        a_smem = warp.load(a_smem, a, (), (0, 0, 0, tile_col), axis=2)
+        a_reg = warp.load(a_reg, a_smem)
+
+        a_reg = warp.map(a_reg, lambda x: x * (1.0 / math.log(2)))
+
+        max_vec_last = warp.copy(max_vec_last.after(tile_col), max_vec)
+        max_vec = warp.row_reduce(max_vec, a_reg, lambda a, b: a.maximum(b))
+        a_reg = warp.map(a_reg, lambda x, idx: (x - max_vec[idx[0], 0, (idx[2]%4)//2]).exp2())
+        max_vec_last = warp.map(max_vec_last, lambda x, idx: (x - max_vec[*idx]).exp2())
+        norm_vec = warp.map(norm_vec, lambda x, idx: x * max_vec_last[*idx])
+        norm_vec = warp.row_reduce(norm_vec, a_reg, lambda a, b: a + b)
+      norm_vec = ker.endrange()
+
+      for tile_col in ker.range(N // BLOCK_SIZE):
+        a_smem = warp.load(a_smem, a, (), (0, 0, 0, tile_col), axis=2)
+        a_reg = warp.load(a_reg, a_smem)
+
+        a_reg = warp.map(a_reg, lambda x: x * (1.0 / math.log(2)))
+        a_reg = warp.map(a_reg, lambda x, idx: (x - max_vec[idx[0], 0, (idx[2]%4)//2]).exp2())
+        a_reg = warp.map(a_reg, lambda x, idx: x / norm_vec[idx[0], 0, (idx[2]%4)//2])
+
+        a_smem = warp.store(a_smem, a_reg)
+        b = warp.store(b, a_smem, (0, 0, 0, tile_col), (), axis=2)
+
+      sink = ker.finish()
+
+    with Context(DEBUG=0):
+      a = Tensor.rand(1, 1, BLOCK_SIZE, N, dtype="float32")
+      b = Tensor.empty(1, 1, BLOCK_SIZE, N, dtype="float32")
+      Tensor.realize(a, b)
+
+    ei = ExecItem(get_runner(Device.DEFAULT, sink), [t.uop.buffer for t in (b, a)])
+    for _ in range(5): ei.run(wait=True)
+    b = b.float()
+
+    ref = a.float().softmax(axis=3)
+
+    np.testing.assert_allclose(b.numpy(), ref.numpy(), atol=1e-5, rtol=1e-5)

 if __name__ == "__main__":
  unittest.main()
@@ -0,0 +1,47 @@
+import ctypes, subprocess, tempfile, unittest
+from tinygrad.helpers import WIN
+from tinygrad.runtime.support.c import Struct
+
+class TestAutogen(unittest.TestCase):
+  def test_packed_struct_sizeof(self):
+    layout = [('a', ctypes.c_char), ('b', ctypes.c_int, 5), ('c', ctypes.c_char)]
+    class X(ctypes.Structure): _fields_, _layout_ = layout, 'gcc-sysv'
+    class Y(ctypes.Structure): _fields_, _pack_, _layout_ = layout, 1, 'ms'
+    class Z(Struct): _packed_, _fields_ = True, layout
+    self.assertNotEqual(ctypes.sizeof(X), 4) # ctypes bug! gcc-13.3.0 says this should have size 4
+    self.assertEqual(ctypes.sizeof(Y), 6)
+    self.assertEqual(ctypes.sizeof(Z), 3)
+    layout = [('a', ctypes.c_int, 31), ('b', ctypes.c_int, 31), ('c', ctypes.c_int, 1), ('d', ctypes.c_int, 1)]
+    class Foo(ctypes.Structure): _fields_, _layout_ = layout, 'gcc-sysv'
+    class Bar(ctypes.Structure): _fields_, _pack_, _layout_ = layout, 1, 'ms'
+    class Baz(Struct): _fields_, _packed_ = layout, True
+    self.assertEqual(ctypes.sizeof(Foo), 12)
+    self.assertEqual(ctypes.sizeof(Bar), 12)
+    self.assertEqual(ctypes.sizeof(Baz), 8)
+
+  @unittest.skipIf(WIN, "doesn't compile on windows")
+  def test_packed_struct_interop(self):
+    class Baz(Struct): pass
+    Baz._packed_ = True
+    Baz._fields_ = [('a', ctypes.c_int, 30), ('b', ctypes.c_int, 30), ('c', ctypes.c_int, 2), ('d', ctypes.c_int, 2)]
+    src = '''
+      struct __attribute__((packed)) baz {
+        int a:30;
+        int b:30;
+        int c:2;
+        int d:2;
+      };
+
+      int test(struct baz x) {
+        return x.a + x.b + x.c + x.d;
+      }
+    '''
+    args = ('-x', 'c', '-fPIC', '-shared')
+    with tempfile.NamedTemporaryFile(suffix=".so") as f:
+      subprocess.check_output(('clang',) + args + ('-', '-o', f.name), input=src.encode('utf-8'))
+      b = Baz(0xAA000, 0x00BB0, 0, 1)
+      test = ctypes.CDLL(f.name).test
+      test.argtypes = [Baz]
+      self.assertEqual(test(b), b.a + b.b + b.c + b.d)
+
+if __name__ == "__main__": unittest.main()
@@ -62,6 +62,7 @@ class TestConv(unittest.TestCase):
    np.testing.assert_allclose(r1.numpy(), np.maximum(out.numpy(), 0), atol=1e-5)
    np.testing.assert_allclose(r2.numpy(), np.where(out.numpy() > 0, out.numpy(), (np.exp(out.numpy()) - 1)), atol=1e-5)

+  @unittest.skip("this test is flaky")
  def test_two_overlapping_binops_no_rerun_wino(self):
    with Context(WINO=1):
      x = Tensor.randn(1,4,16,16)
@@ -81,20 +81,20 @@ class TestCompiler(unittest.TestCase):
  def test_compile_cached(self):
    diskcache_put("key", "123", None) # clear cache
    getenv.cache_clear()
-    with Context(DISABLE_COMPILER_CACHE=0):
+    with Context(CCACHE=1):
      self.assertEqual(MockCompiler("key").compile_cached("123"), str.encode("123"))
      self.assertEqual(diskcache_get("key", "123"), str.encode("123"))

  def test_compile_cached_disabled(self):
    diskcache_put("disabled_key", "123", None) # clear cache
    getenv.cache_clear()
-    with Context(DISABLE_COMPILER_CACHE=1):
+    with Context(CCACHE=0):
      self.assertEqual(MockCompiler("disabled_key").compile_cached("123"), str.encode("123"))
      self.assertIsNone(diskcache_get("disabled_key", "123"))

  def test_device_compile(self):
    getenv.cache_clear()
-    with Context(DISABLE_COMPILER_CACHE=1):
+    with Context(CCACHE=0):
      a = Tensor([0.,1.], device=Device.DEFAULT).realize()
      (a + 1).realize()

@@ -211,5 +211,39 @@ class TestPatternMatcher(unittest.TestCase):
      return u.src[0]
    for a,b in zip(simple_src(a), simple_src(b)): self._assert_eq_upat(a, b)

+class TestAlgebraic(unittest.TestCase):
+  def test_plus_0(self):
+    pm = PatternMatcher([
+      (UPat.var("x") + 0, UPat.var("x")), # x+0 -> x
+    ])
+    expr = UOp.const(dtypes.int, 4)+0
+    print(expr)
+    self.assertEqual(pm.rewrite(expr), UOp.const(dtypes.int, 4))
+
+  def test_div_mul(self):
+    pm = PatternMatcher([
+      ((UPat.var("x") * UPat.var("x2")) / UPat.var("x2"), UPat.var("x")), # (x*x2)/x2 -> x
+    ])
+    expr = UOp.const(dtypes.float, 4)/2*2
+    print(expr)
+    self.assertEqual(pm.rewrite(expr), UOp.const(dtypes.int, 4))
+
+  def test_mul_is_and(self):
+    pm = PatternMatcher([
+      (UPat.var('x', dtype=dtypes.bool) * UPat.var('y', dtype=dtypes.bool), UPat.var('x') & UPat.var('y')),
+    ])
+    expr = UOp.const(dtypes.bool, True)*UOp.const(dtypes.bool, True)
+    print(expr)
+    self.assertEqual(pm.rewrite(expr), UOp.const(dtypes.bool, True)&UOp.const(dtypes.bool, True))
+
+  def test_div_neg_1(self):
+    pm = PatternMatcher([
+      (UPat.var("x") // -1, UPat.var("x") * -1), # x//-1 -> x * -1
+    ])
+    expr = UOp.const(dtypes.float, 4)//-1
+    print(expr)
+    self.assertEqual(pm.rewrite(expr), UOp.const(dtypes.int, 4) * -1)
+
+
 if __name__ == '__main__':
  unittest.main(verbosity=2)
@@ -366,8 +366,8 @@ def load_profile(lst:list[ProfileEvent]) -> dict:
        else: v["events"].append({"event":"free", "ts":ts, "key":key, "arg": {"users":[u("<IIBB") for _ in range(u("<I")[0])]}})
  return {"dur":total_dur, "peak":global_peak, "layout":layout, "markers":markers}

-class TestVizProfiler(unittest.TestCase):
-  def test_perfetto_node(self):
+class TestVizProfiler(BaseTestViz):
+  def test_node(self):
    prof = [ProfileRangeEvent(device='NV', name='E_2', st=decimal.Decimal(1000), en=decimal.Decimal(1010), is_copy=False),
            ProfileDeviceEvent(device='NV', comp_tdiff=decimal.Decimal(-1000), copy_tdiff=decimal.Decimal(-100))]

@@ -381,7 +381,7 @@ class TestVizProfiler(unittest.TestCase):
    self.assertEqual(event['dur'], 10)
    assert event['ref'] is None

-  def test_perfetto_copy_node(self):
+  def test_copy_node(self):
    prof = [ProfileRangeEvent(device='NV', name='COPYxx', st=decimal.Decimal(1000), en=decimal.Decimal(1010), is_copy=True),
            ProfileRangeEvent(device='NV:2', name='COPYxx', st=decimal.Decimal(1000), en=decimal.Decimal(1010), is_copy=True),
            ProfileDeviceEvent(device='NV', comp_tdiff=decimal.Decimal(-1000), copy_tdiff=decimal.Decimal(-100)),
@@ -399,7 +399,7 @@ class TestVizProfiler(unittest.TestCase):

    self.assertEqual(j["dur"], (event2["st"]+event2["dur"])-event["st"])

-  def test_perfetto_graph(self):
+  def test_graph(self):
    prof = [ProfileDeviceEvent(device='NV', comp_tdiff=decimal.Decimal(-1000), copy_tdiff=decimal.Decimal(-100)),
            ProfileDeviceEvent(device='NV:1', comp_tdiff=decimal.Decimal(-500), copy_tdiff=decimal.Decimal(-50)),
            ProfileGraphEvent(ents=[ProfileGraphEntry(device='NV', name='E_25_4n2', st_id=0, en_id=1, is_copy=False),
@@ -436,6 +436,12 @@ class TestVizProfiler(unittest.TestCase):
    sz = len(get_profile(prof))
    self.assertLessEqual(sz/n_events, 26)

+  def test_calltrace(self):
+    def fxn(): return Tensor.empty(10).mul(2).realize()
+    fxn()
+    trace = get_viz_list()[0]["steps"][0]["trace"]
+    assert any(fxn.__code__.co_filename == f and fxn.__code__.co_firstlineno == l for f,l,*_ in trace), str(trace)
+
  # can pack up to 1hr 11 min of trace events
  def test_trace_duration(self):
    dur_mins = 72
@@ -81,10 +81,10 @@ def hand_coded_optimizations(k:Scheduler) -> Scheduler:
            return k

  # are we grouping? (requires local shape support)
-  if resolve(prod(k.output_shape[i] for i in k.upcastable_dims) <= (128 if NOLOCALS else 2048), False):
-    for sz in [16]:
+  if resolve(prod(k.output_shape[i] for i in k.upcastable_dims) <= (240 if NOLOCALS else 2048), False):
+    for axis, sz in itertools.product((0, 1, 2), (16,)):
      try:
-        k.apply_opt(Opt(OptOps.GROUPTOP, 0, sz))
+        k.apply_opt(Opt(OptOps.GROUPTOP, axis, sz))
        break
      except KernelOptError: pass

@@ -2,8 +2,7 @@ from __future__ import annotations
 import math, itertools
 from collections import defaultdict
 from typing import cast, Final
-from tinygrad.uop.ops import PatternMatcher, UPat, Ops, UOp, KernelInfo, graph_rewrite, AxisType, ssimplify, GroupOp
-from tinygrad.uop.ops import axis_letters, axis_colors, axis_to_pos
+from tinygrad.uop.ops import PatternMatcher, UPat, Ops, UOp, KernelInfo, graph_rewrite, AxisType, ssimplify, GroupOp, axis_letters, axis_colors
 from tinygrad.device import Buffer
 from tinygrad.dtype import dtypes, ImageDType
 from tinygrad.helpers import colored, BEAM, getenv, DEBUG, to_function_name, NOOPT, argsort, round_up, prod, merge_dicts, get_single_element, flatten
@@ -13,6 +12,10 @@ from tinygrad.renderer import Renderer

 remove_tags = PatternMatcher([(UPat(GroupOp.All, name="x"), lambda x: x.replace(tag=None) if x.tag is not None else None)])

+# NOTE: LOCAL and GROUP_REDUCE have the same priority. the order here matters
+axis_to_pos = {AxisType.LOOP: -1, AxisType.THREAD: 0, AxisType.GLOBAL: 0, AxisType.WARP: 1, AxisType.LOCAL: 2, AxisType.UPCAST: 3,
+               AxisType.GROUP_REDUCE: 2, AxisType.REDUCE: 4, AxisType.UNROLL: 5}
+
 class Scheduler:
  def __init__(self, ast:UOp, ren:Renderer):
    self.ast, self.ren = ast, ren
@@ -4,7 +4,7 @@ from collections import defaultdict
 from typing import Any, Generic, TypeVar, Iterator, Sequence, cast, Generator
 import importlib, inspect, functools, pathlib, os, platform, contextlib, sys, re, atexit, pickle, decimal
 from tinygrad.helpers import CI, OSX, LRU, getenv, diskcache_get, diskcache_put, DEBUG, GlobalCounters, flat_mv, PROFILE, temp, colored, CPU_LLVM
-from tinygrad.helpers import Context, DISABLE_COMPILER_CACHE, ALLOW_DEVICE_USAGE, MAX_BUFFER_SIZE, cpu_events, ProfileEvent, ProfilePointEvent, dedup
+from tinygrad.helpers import Context, CCACHE, ALLOW_DEVICE_USAGE, MAX_BUFFER_SIZE, cpu_events, ProfileEvent, ProfilePointEvent, dedup
 from tinygrad.helpers import unwrap_class_type, suppress_finalizing, AMD_LLVM, select_first_inited
 from tinygrad.dtype import DType, ImageDType, PtrDType, dtypes, _to_np_dtype
 from tinygrad.renderer import Renderer
@@ -266,7 +266,7 @@ class LRUAllocator(Allocator, Generic[DeviceType]):
 class CompileError(Exception): pass

 class Compiler:
-  def __init__(self, cachekey:str|None=None): self.cachekey = None if DISABLE_COMPILER_CACHE else cachekey
+  def __init__(self, cachekey:str|None=None): self.cachekey = cachekey if CCACHE else None
  def compile(self, src:str) -> bytes: return src.encode()   # NOTE: empty compiler is the default
  def compile_cached(self, src:str) -> bytes:
    if self.cachekey is None or (lib := diskcache_get(self.cachekey, src)) is None:
@@ -3,7 +3,7 @@ import time, pprint, random, itertools, math
 from dataclasses import dataclass, replace, field
 from tinygrad.helpers import all_same, colored, DEBUG, GlobalCounters, ansilen, BEAM, NOOPT, all_int, CAPTURING, Metadata, TRACEMETA, TracingKey
 from tinygrad.helpers import DEVECTORIZE, time_to_str, VALIDATE_WITH_CPU, getenv, cpu_profile, PROFILE, ProfilePointEvent, cpu_events, prod, Context
-from tinygrad.helpers import unwrap
+from tinygrad.helpers import unwrap, disable_gc
 from tinygrad.uop.ops import Ops, PatternMatcher, UOp, UPat, sym_infer, graph_rewrite, print_uops, track_rewrites, KernelInfo, pyrender
 from tinygrad.device import Device, Buffer
 from tinygrad.renderer import Renderer, ProgramSpec, Estimates
@@ -13,6 +13,7 @@ from tinygrad.codegen.opt import Opt

 # **************** Program Creation ****************

+@disable_gc()
@track_rewrites(name=lambda *args,ret,**kwargs: TracingKey(ret.name, (ret.function_name, ret.ast), ret=ret), replay=True)
 def get_program(ast:UOp, renderer:Renderer|None=None, opts:list[Opt]|None=None) -> ProgramSpec:
  """
@@ -1,5 +1,5 @@
 from __future__ import annotations
-import os, functools, platform, time, re, contextlib, operator, hashlib, pickle, sqlite3, tempfile, pathlib, string, ctypes, sys, gzip, getpass
+import os, functools, platform, time, re, contextlib, operator, hashlib, pickle, sqlite3, tempfile, pathlib, string, ctypes, sys, gzip, getpass, gc
 import urllib.request, subprocess, shutil, math, types, copyreg, inspect, importlib, decimal, itertools
 from dataclasses import dataclass, field
 from typing import ClassVar, Iterable, Any, TypeVar, Callable, Sequence, TypeGuard, Iterator, Generic, Generator, cast, overload
@@ -173,13 +173,12 @@ TRANSCENDENTAL, NOLOCALS = ContextVar("TRANSCENDENTAL", 1), ContextVar("NOLOCALS
 SPLIT_REDUCEOP, NO_MEMORY_PLANNER, RING = ContextVar("SPLIT_REDUCEOP", 1), ContextVar("NO_MEMORY_PLANNER", 0), ContextVar("RING", 1)
 PICKLE_BUFFERS, LRU = ContextVar("PICKLE_BUFFERS", 1), ContextVar("LRU", 1)
 CACHELEVEL, IGNORE_BEAM_CACHE, DEVECTORIZE = ContextVar("CACHELEVEL", 2), ContextVar("IGNORE_BEAM_CACHE", 0), ContextVar("DEVECTORIZE", 1)
-DISABLE_COMPILER_CACHE = ContextVar("DISABLE_COMPILER_CACHE", 0)
 VALIDATE_WITH_CPU, DISABLE_FAST_IDIV = ContextVar("VALIDATE_WITH_CPU", 0), ContextVar("DISABLE_FAST_IDIV", 0)
 CORRECT_DIVMOD_FOLDING, FUSE_OPTIM = ContextVar("CORRECT_DIVMOD_FOLDING", 0), ContextVar("FUSE_OPTIM", 0)
 ALLOW_DEVICE_USAGE, MAX_BUFFER_SIZE = ContextVar("ALLOW_DEVICE_USAGE", 1), ContextVar("MAX_BUFFER_SIZE", 0)
 EMULATE = ContextVar("EMULATE", "")
 CPU_COUNT = ContextVar("CPU_COUNT", max(1, len(os.sched_getaffinity(0)) if hasattr(os, "sched_getaffinity") else (os.cpu_count() or 1)))
-CPU_LLVM, CPU_LVP, AMD_LLVM = ContextVar("CPU_LLVM", 0), ContextVar("CPU_LVP", 0), ContextVar("AMD_LLVM", 1)
+CPU_LLVM, CPU_LVP, AMD_LLVM = ContextVar("CPU_LLVM", 0), ContextVar("CPU_LVP", 0), ContextVar("AMD_LLVM", 0)
 VIZ = PROFILE = ContextVar("VIZ", 0)
 SPEC = ContextVar("SPEC", 1)
 # TODO: disable by default due to speed
@@ -188,6 +187,8 @@ PCONTIG = ContextVar("PCONTIG", 0)  # partial contiguous in rangeify
 DEBUG_RANGEIFY = ContextVar("DEBUG_RANGEIFY", 0)
 # set to 1, this uses tuplize in the linearizer sort order
 TUPLE_ORDER = ContextVar("TUPLE_ORDER", 1)
+# set to 0 to disable the compiler cache
+CCACHE = ContextVar("CCACHE", 1)

@dataclass(frozen=True)
 class Metadata:
@@ -240,11 +241,29 @@ class Profiling(contextlib.ContextDecorator):

 def perf_counter_us() -> decimal.Decimal: return decimal.Decimal(time.perf_counter_ns())/1000

+@functools.cache
+def lines(fn) -> list[str]:
+  try:
+    with open(fn, encoding="utf-8") as f: return f.readlines()
+  except (FileNotFoundError, OSError): return []
+
+def printable(loc:tuple[str, int]) -> str:
+  try: return lines(loc[0])[loc[1]-1].strip()
+  except IndexError: return "<missing>"
+
+def get_stacktrace(frm, max_frames=30) -> tuple[tuple, ...]:
+  ret:list[tuple] = []
+  for i in range(max_frames):
+    if (frm:=frm.f_back) is None: break
+    ret.append(((fc:=frm.f_code).co_filename, frm.f_lineno, fc.co_name, printable((fc.co_filename, frm.f_lineno))))
+  return tuple(ret)
+
@dataclass(frozen=True)
 class TracingKey:
  display_name:str                       # display name of this trace event
  keys:tuple[Any, ...]=()                # optional keys to search for related traces
  ret:Any=None
+  tb:tuple[tuple, ...]|None=field(default_factory=lambda: get_stacktrace(sys._getframe(1)) if VIZ else None)

 class ProfileEvent: pass

@@ -361,10 +380,12 @@ def fetch(url:str, name:pathlib.Path|str|None=None, subdir:str|None=None, gunzip

 # *** Exec helpers

+def system(cmd, **kwargs): return subprocess.check_output(cmd.split(), **kwargs).decode().strip()
+
 def cpu_objdump(lib, objdump_tool='objdump'):
  with tempfile.NamedTemporaryFile(delete=True) as f:
    pathlib.Path(f.name).write_bytes(lib)
-    print(subprocess.check_output([objdump_tool, '-d', f.name]).decode('utf-8'))
+    print(system(f"{objdump_tool} -d {f.name}"))

 def capstone_flatdump(lib: bytes):
  try: import capstone
@@ -442,6 +463,13 @@ class tqdm(Generic[T]):
 class trange(tqdm):
  def __init__(self, n:int, **kwargs): super().__init__(iterable=range(n), total=n, **kwargs)

+class disable_gc(contextlib.ContextDecorator):
+  def __enter__(self):
+    self._was_enabled = gc.isenabled()
+    if self._was_enabled: gc.disable()
+  def __exit__(self, *exc):
+    if self._was_enabled: gc.enable()
+
 # *** universal support for code object pickling

 def _reconstruct_code(*args): return types.CodeType(*args)
@@ -2,7 +2,8 @@
 import functools
 from typing import TypeAlias, TYPE_CHECKING, Self
 from tinygrad.uop import Ops
-from tinygrad.helpers import prod, argfix, flatten, dedup
+from tinygrad.helpers import prod, argfix, flatten, dedup, make_tuple, ceildiv
+from tinygrad.uop.ops import resolve, smax
 if TYPE_CHECKING: from tinygrad.uop.ops import UOp
 sint: TypeAlias = "UOp | int"

@@ -326,3 +327,24 @@ class MovementMixin:
    expanded_shape = flatten([[s] if r == 1 else [r, s] for r,s in zip(repeats, base_shape)])
    final_shape = [r*s for r,s in zip(repeats, base_shape)]
    return self.reshape(unsqueezed_shape).expand(expanded_shape).reshape(final_shape)
+
+  # **** pool level ****
+
+  def _pool(self, k_:tuple[sint, ...], stride:int|tuple[int, ...]=1, dilation:int|tuple[int, ...]=1) -> Self:
+    assert len(self.shape) >= len(k_), f"can't pool {self.shape} with {k_}"
+    s_, d_ = make_tuple(stride, len(k_)), make_tuple(dilation, len(k_))
+    assert len(k_) == len(s_) == len(d_), f"stride/dilation mismatch kernel:{k_} stride:{s_} dilation:{d_}"
+    noop, i_ = [None] * (self.ndim-len(k_)), self.shape[-len(k_):]
+    assert all(resolve(d*(k-1)+1 <= i) for k,d,i in zip(k_,d_,i_)), "kernel size cannot be greater than actual input size"
+    o_ = [ceildiv(i-d*(k-1), s) for i,d,k,s in zip(i_,d_,k_,s_)]
+    # input size scaling factor to make sure shrink for stride is possible
+    f_ = [smax(1, ceildiv(o*s - d, i)) for o,s,i,d in zip(o_,s_,i_,d_)]
+    # repeats such that we don't need padding
+    x = self.repeat([1]*len(noop) + [ceildiv(k*(i*f+d),i) for k,i,d,f in zip(k_,i_,d_,f_)])
+    # handle dilation
+    x = x.shrink_to(noop + [k*(i*f+d) for k,i,d,f in zip(k_,i_,d_,f_)]).reshape(noop + flatten((k,(i*f+d)) for k,i,d,f in zip(k_,i_,d_,f_)))
+    # handle stride
+    x = x.shrink_to(noop + flatten((k,o*s) for k,o,s in zip(k_,o_,s_))).reshape(noop + flatten((k,o,s) for k,o,s in zip(k_,o_,s_)))
+    x = x.shrink_to(noop + flatten((k,o,1) for k,o in zip(k_,o_))).reshape(noop + flatten((k,o) for k,o in zip(k_,o_)))
+    # permute to move reduce to the end
+    return x.permute(*range(len(noop)), *[len(noop)+i*2+1 for i in range(len(i_))], *[len(noop)+i*2 for i in range(len(i_))])
@@ -1124,6 +1124,16 @@ def get_onnx_ops() -> dict[str, types.FunctionType|dict[OpSetId, types.FunctionT
    return output.flatten(start_dim=2) if len(original_input_shape) == 3 else output.permute(0, 2, 1, 3)

  # ***** Indexing Ops *****
+  def NonZero(x:Tensor):
+    mask = (x!=0).flatten()
+    flat_idx = Tensor.arange(mask.numel(), dtype=dtypes.int64, device=x.device).masked_select(mask)
+    if flat_idx.ndim == 0: flat_idx = flat_idx.reshape(1)
+    if x.ndim == 0:
+      return Tensor.zeros((0, flat_idx.shape[0]), dtype=dtypes.int64, device=x.device, requires_grad=False)
+    strides = [prod(int(s) for s in x.shape[i+1:]) if i+1 < x.ndim else 1 for i in range(x.ndim)]
+    coords = [((flat_idx // stride) % int(dim)) for stride, dim in zip(strides, x.shape)]
+    return Tensor.stack(*coords, dim=0)
+
  def ArrayFeatureExtractor(x:Tensor, indices:Tensor): return x[..., indices]

  def Gather(x:Tensor, indices:Tensor, axis:int=0):
@@ -450,16 +450,9 @@ class AMDRenderer(CStyleLanguage):
      ]) + base_rewrite
  def __reduce__(self): return self.__class__, (self.arch,)

-  # language options
-  ockl = [(f"__ockl_get_{name}", "unsigned int", "size_t", "const") for name in ["local_id", "group_id", "local_size"]]
-  ocml = [(f"__ocml_{name}_f{n}", f"{dt}, {dt}" if "fmax" == name else dt, dt, atr)
-            for dt, n in [(dtype.name, dtype.itemsize * 8) for dtype in [dtypes.float, dtypes.double, dtypes.half]]
-            for name, atr in [("fmax", "const"), ("exp2", "pure"), ("log2", "pure"), ("sqrt", "const"), ("sin", ""), ("trunc", "")]]
-
-  kernel_typedef = "\n".join(f'extern "C" __attribute__((device{f", {atr}" if atr else ""})) {dto} {meth}({dti});' for meth,dti,dto,atr in ockl+ocml)
  # https://clang.llvm.org/docs/AttributeReference.html#amdgpu-flat-work-group-size
  # NOTE: this makes hlb_cifar10 twice as fast, there may be more gains in tweaking these parameters
-  kernel_typedef += '\nextern "C" __attribute__((global)) void __attribute__((amdgpu_flat_work_group_size(1, {launch_bounds})))'
+  kernel_typedef = 'extern "C" __attribute__((global)) void __attribute__((amdgpu_flat_work_group_size(1, {launch_bounds})))'
  code_for_workitem = {"g": lambda x: f"__ockl_get_group_id({x})", "l": lambda x: f"__ockl_get_local_id({x})",
                       "i": lambda x: f"(__ockl_get_group_id({x})*__ockl_get_local_size({x})+__ockl_get_local_id({x}))"}
  code_for_op = { **CStyleLanguage.code_for_op,
@@ -490,15 +483,25 @@ class AMDRenderer(CStyleLanguage):
           f"{vec} make_{vec}({', '.join([f'{scal} {x}' for x in _nms[:dtype.count]])}) {{ return {{ {', '.join(_nms[:dtype.count])} }}; }}"

  def render_kernel(self, function_name, kernel, bufs, uops, prefix=None) -> str:
-    prefix = ["#define INFINITY (__builtin_inff())","#define NAN (__builtin_nanf(\"\"))","typedef long unsigned int size_t;","#define half _Float16"]
+    prefix, ockl = [], []
    type_map = { dtypes.bfloat16: "bf16", dtypes.float: "f32", dtypes.half: "f16", dtypes.fp8e4m3: "_fp8_fp8", dtypes.fp8e5m2: "_bf8_bf8" }
    used_dtypes = uops_to_dtypes(uops)
+    if any(u.op is Ops.CONST and not math.isfinite(u.arg) for u in uops):
+      prefix += ["#define INFINITY (__builtin_inff())", "#define NAN (__builtin_nanf(\"\"))"]
+    if any(u.op is Ops.SPECIAL for u in uops):
+      prefix.append("typedef long unsigned int size_t;")
+      ockl = [(f"__ockl_get_{name}", "unsigned int", "size_t", "const") for name in ["local_id", "group_id", "local_size"]]
+    ocml_ops = {Ops.EXP2: ("exp2", "pure"), Ops.LOG2: ("log2", "pure"), Ops.SQRT: ("sqrt", "const"), Ops.SIN: ("sin", ""), Ops.TRUNC: ("trunc", "")}
+    ocml = [(f"__ocml_{ocml_ops[op][0]}_f{dt.itemsize * 8}", dt.name, dt.name, ocml_ops[op][1])
+      for op, dt in dedup((u.op, u.dtype.scalar()) for u in uops) if op in ocml_ops and dt in (dtypes.half, dtypes.float, dtypes.double)]
    if any(dt.scalar() == dtypes.bfloat16 for dt in used_dtypes): prefix.append("typedef unsigned short hip_bfloat16;")
+    if any(dt.scalar() == dtypes.half for dt in used_dtypes): prefix.append("#define half _Float16")
    if any(dt.scalar() in dtypes.fp8s for dt in used_dtypes):
      prefix += ["typedef unsigned char hip_bf8;", "typedef unsigned char hip_fp8;"]
      prefix.append("""static inline __attribute__((device)) unsigned char f32_to_fp8(float v, int is_bf8) {
  v = (((*(unsigned*)&v)&0x7F800000)!=0x7F800000)?__builtin_amdgcn_fmed3f(v,is_bf8?57344.0f:448.0f,is_bf8?-57344.0f:-448.0f) : v;
  return (unsigned char)(is_bf8?__builtin_amdgcn_cvt_pk_bf8_f32(v,v,0,false):__builtin_amdgcn_cvt_pk_fp8_f32(v,v,0,false));\n}""")
+    prefix += [f'extern "C" __attribute__((device{f", {atr}" if atr else ""})) {dto} {meth}({dti});' for meth,dti,dto,atr in ockl+ocml]
    prefix += [self.render_vector_prefix(dt) for dt in used_dtypes if dt.count > 1]

    for name, (N, M, K), dtype_in, dtype_out, _, _, _, _ in wmma_args(uops): # TODO: handle TCs f32_bf16 and bf16_bf16 w/ wrapper
@@ -0,0 +1,17 @@
+import importlib, pathlib
+from tinygrad.helpers import system
+
+root = (here:=pathlib.Path(__file__).parent).parents[2]
+
+def load(name, dll, files, **kwargs):
+  if not (f:=(root/(path:=kwargs.pop("path", __name__)).replace('.','/')/f"{name}.py")).exists():
+    files = files() if callable(files) else files
+    f.write_text(importlib.import_module("tinygrad.runtime.support.autogen").gen(dll, files, **kwargs))
+  return importlib.import_module(f"{path}.{name.replace('/', '.')}")
+
+def __getattr__(nm):
+  match nm:
+    case "libc": return load("libc", ["find_library('c')"], lambda: (
+      [i for i in system("dpkg -L libc6-dev").split() if 'sys/mman.h' in i or 'sys/syscall.h' in i] +
+      ["/usr/include/string.h", "/usr/include/elf.h", "/usr/include/unistd.h", "/usr/include/asm-generic/mman-common.h"]), use_errno=True)
+    case _: raise AttributeError(f"no such autogen: {nm}")
@@ -594,7 +594,8 @@ class AMDProgram(HCQProgram):
      self.dev.synchronize()

      for se, buf in enumerate(self.dev.sqtt_buffers):
-        wptr = ((self.dev.sqtt_wptrs.cpu_view().view(fmt='I')[se]&0x1FFFFFFF)-(((buf.va_addr//32)&0x1FFFFFFF) if self.dev.target[0] == 11 else 0))*32
+        wptr = (self.dev.sqtt_wptrs.cpu_view().view(fmt='I')[se] & 0x1FFFFFFF) * 32
+        if self.dev.target[:2] == (11, 0): wptr -= ((buf.va_addr // 32) & 0x1FFFFFFF) * 32

        if DEBUG >= 5: print(f'\t{self.dev.device}: SE {se} blob size {wptr:#x}')
        assert wptr >= 0 and wptr <= buf.size, f"{wptr} > {buf.size}, should never happen"
@@ -785,7 +786,7 @@ class PCIIface(PCIIfaceBase):
  gpus:ClassVar[list[str]] = []

  def __init__(self, dev, dev_id):
-    super().__init__(dev, dev_id, vendor=0x1002, devices=[0x744c, 0x7480, 0x7550, 0x7590], bars=[0, 2, 5], vram_bar=0,
+    super().__init__(dev, dev_id, vendor=0x1002, devices=[(0xffff, [0x744c, 0x7480, 0x7550, 0x7590])], bars=[0, 2, 5], vram_bar=0,
      va_start=AMMemoryManager.va_allocator.base, va_size=AMMemoryManager.va_allocator.size)
    self._setup_adev(self.pci_dev)
    self.pci_dev.write_config(pci.PCI_COMMAND, self.pci_dev.read_config(pci.PCI_COMMAND, 2) | pci.PCI_COMMAND_MASTER, 2)
@@ -911,8 +912,8 @@ class AMDDevice(HCQCompiled):
    max_copy_size = 0x40000000 if self.iface.ip_versions[am.SDMA0_HWIP][0] >= 5 else 0x400000
    self.sdma_queue = self.create_queue(kfd.KFD_IOC_QUEUE_TYPE_SDMA, 0x200 if self.is_usb() else (16 << 20))

-    compilers:list[CompilerPairT] = [(functools.partial(AMDLLVMRenderer, self.arch), functools.partial(AMDLLVMCompiler, self.arch)),
-                                     (functools.partial(AMDRenderer, self.arch), functools.partial(HIPCompiler, self.arch))]
+    compilers:list[CompilerPairT] = [(functools.partial(AMDRenderer, self.arch), functools.partial(HIPCompiler, self.arch)),
+                                     (functools.partial(AMDLLVMRenderer, self.arch), functools.partial(AMDLLVMCompiler, self.arch))]

    super().__init__(device, AMDAllocator(self), compilers, functools.partial(AMDProgram, self), AMDSignal,
                     functools.partial(AMDComputeAQLQueue if self.is_aql else AMDComputeQueue, self),
@@ -5,7 +5,7 @@ from tinygrad.device import BufferSpec, Compiled, Allocator, Compiler
 from tinygrad.runtime.ops_cpu import CPUAllocator
 from tinygrad.dtype import dtypes, DType, PtrDType
 from tinygrad.uop.ops import Ops, UOp
-from tinygrad.helpers import getenv, round_up, mv_address, to_mv, cpu_objdump, DEBUG
+from tinygrad.helpers import getenv, round_up, mv_address, to_mv, cpu_objdump, system, DEBUG
 from tinygrad.renderer.cstyle import ClangRenderer
 from tinygrad.runtime.autogen import libc, qcom_dsp
 if getenv("IOCTL"): import extra.dsp.run # noqa: F401 # pylint: disable=unused-import
@@ -123,10 +123,9 @@ class ClangCompiler(Compiler):

  def compile(self, src:str) -> bytes:
    # TODO: remove file write. sadly clang doesn't like the use of /dev/stdout here
-    with tempfile.NamedTemporaryFile(delete=True) as output_file:
-      subprocess.check_output([getenv("CC", 'clang'), *self.args, '-O2', '-Wall', '-Werror', '-x', 'c', '-fPIC', '-ffreestanding', '-nostdlib',
-                               '-', '-o', str(output_file.name)], input=src.encode('utf-8'))
-      return pathlib.Path(output_file.name).read_bytes()
+    with tempfile.NamedTemporaryFile(delete=True) as f:
+      system(f"{getenv('CC','clang')} {' '.join(self.args)} -O2 -Wall -Werror -x c -fPIC -ffreestanding -nostdlib - -o {f.name}", input=src.encode())
+      return pathlib.Path(f.name).read_bytes()

  def disassemble(self, lib:bytes): return cpu_objdump(lib, self.objdump_tool)

@@ -456,8 +456,8 @@ class PCIIface(PCIIfaceBase):
  gpus:ClassVar[list[str]] = []

  def __init__(self, dev, dev_id):
-    super().__init__(dev, dev_id, vendor=0x10de, devices=[0x2204, 0x2684, 0x2b85], bars=[0, 1], vram_bar=1,
-      va_start=NVMemoryManager.va_allocator.base, va_size=NVMemoryManager.va_allocator.size)
+    super().__init__(dev, dev_id, vendor=0x10de, devices=[(0xff00, [0x2200, 0x2400, 0x2500, 0x2600, 0x2700, 0x2800, 0x2b00, 0x2c00, 0x2d00, 0x2f00])],
+      bars=[0, 1], vram_bar=1, va_start=NVMemoryManager.va_allocator.base, va_size=NVMemoryManager.va_allocator.size)
    if not OSX: System.reserve_hugepages(64)

    self.pci_dev.write_config(pci.PCI_COMMAND, self.pci_dev.read_config(pci.PCI_COMMAND, 2) | pci.PCI_COMMAND_MASTER, 2)
@@ -1,5 +1,5 @@
 from __future__ import annotations
-import os, ctypes, functools, mmap, struct, array, math, sys, weakref
+import os, ctypes, functools, mmap, struct, array, math, sys, weakref, contextlib
 assert sys.platform != 'win32'
 from types import SimpleNamespace
 from typing import Any, cast
@@ -9,16 +9,18 @@ from tinygrad.runtime.support.hcq import FileIOInterface, MMIOInterface
 from tinygrad.runtime.autogen import kgsl, adreno
 from tinygrad.runtime.ops_cl import CLCompiler, CLDevice
 from tinygrad.renderer.cstyle import QCOMRenderer
-from tinygrad.helpers import getenv, mv_address, to_mv, round_up, data64_le, prod, fromimport, cpu_profile
+from tinygrad.helpers import getenv, mv_address, to_mv, round_up, data64_le, prod, fromimport, cpu_profile, lo32, PROFILE
+from tinygrad.runtime.support.system import System
 if getenv("IOCTL"): import extra.qcom_gpu_driver.opencl_ioctl  # noqa: F401  # pylint: disable=unused-import

 BUFTYPE_BUF, BUFTYPE_TEX, BUFTYPE_IBO = 0, 1, 2

 #Parse C-style defines: <regname>_<field_x>__SHIFT and <regname>_<field_y>__MASK from the adreno module into the following format:
 # qreg.<regname>(<field_x>=..., <field_y>=..., ..., <field_n>=...)
-def _qreg_exec(reg, __val=0, **kwargs):
+def _qreg_exec(__reg, __val=0, **kwargs):
  for k, v in kwargs.items():
-    __val |= (getattr(adreno, f'{reg[4:]}_{k.upper()}') if v else 0) if type(v) is bool else (v << getattr(adreno, f'{reg[4:]}_{k.upper()}__SHIFT'))
+    reg_name = f"{__reg[4:]}_{k.removeprefix('_').upper()}"
+    __val |= (getattr(adreno, reg_name) if v else 0) if type(v) is bool else (v << getattr(adreno, f'{reg_name}__SHIFT'))
  return __val
 qreg: Any = type("QREG", (object,), {name[4:].lower(): functools.partial(_qreg_exec, name) for name in adreno.__dict__.keys() if name[:4] == 'REG_'})

@@ -67,18 +69,20 @@ class QCOMComputeQueue(HWQueue):
    self._cache_flush(write_back=True, invalidate=True, sync=True, memsync=True)
    return self

-  def signal(self, signal:QCOMSignal, value=0, ts=False):
+  def signal(self, signal:QCOMSignal, value=0):
    self.cmd(adreno.CP_WAIT_FOR_IDLE)
    if self.dev.gpu_id[:2] < (7, 3):
-      self.cmd(adreno.CP_EVENT_WRITE, qreg.cp_event_write_0(event=adreno.CACHE_FLUSH_TS, timestamp=ts),
-               *data64_le(signal.timestamp_addr if ts else signal.value_addr), qreg.cp_event_write_3(value & 0xFFFFFFFF))
+      self.cmd(adreno.CP_EVENT_WRITE, qreg.cp_event_write_0(event=adreno.CACHE_FLUSH_TS), *data64_le(signal.value_addr), lo32(value))
      self._cache_flush(write_back=True, invalidate=False, sync=False, memsync=False)
    else:
      # TODO: support devices starting with 8 Gen 1. Also, 700th series have convenient CP_GLOBAL_TIMESTAMP and CP_LOCAL_TIMESTAMP
      raise RuntimeError('CP_EVENT_WRITE7 is not supported')
    return self

-  def timestamp(self, signal:QCOMSignal): return self.signal(signal, 0, ts=True)
+  def timestamp(self, signal:QCOMSignal):
+    self.cmd(adreno.CP_WAIT_FOR_IDLE)
+    self.cmd(adreno.CP_REG_TO_MEM, qreg.cp_reg_to_mem_0(reg=adreno.REG_A6XX_CP_ALWAYS_ON_COUNTER, cnt=2, _64b=True),*data64_le(signal.timestamp_addr))
+    return self

  def wait(self, signal:QCOMSignal, value=0):
    self.cmd(adreno.CP_WAIT_REG_MEM, qreg.cp_wait_reg_mem_0(function=adreno.WRITE_GE, poll=adreno.POLL_MEMORY),*data64_le(signal.value_addr),
@@ -345,6 +349,9 @@ class QCOMDevice(HCQCompiled):
    # a7xx start with 730x or 'Cxxx', a8xx starts 'Exxx'
    if self.gpu_id[:2] >= (7, 3): raise RuntimeError(f"Unsupported GPU: chip_id={info.chip_id:#x}")

+    if PROFILE and self.gpu_id[:2] < (7, 3):
+      System.write_sysfs("/sys/class/kgsl/kgsl-3d0/idle_timer", value="4000000000", msg="Failed to disable suspend mode", expected="4294967276")
+
    compilers = [(QCOMRenderer, functools.partial(QCOMCompiler, device))]
    super().__init__(device, QCOMAllocator(self), compilers, functools.partial(QCOMProgram, self), QCOMSignal,
                     functools.partial(QCOMComputeQueue, self), None)
@@ -369,3 +376,7 @@ class QCOMDevice(HCQCompiled):
      self.synchronize()
      self._gpu_free(self._stack)
      self._stack = self._gpu_alloc(sz)
+
+  def _at_profile_finalize(self):
+    super()._at_profile_finalize()
+    with contextlib.suppress(RuntimeError): System.write_sysfs("/sys/class/kgsl/kgsl-3d0/idle_timer", "10", "Failed to reenable suspend mode")
@@ -0,0 +1,126 @@
+import ctypes.util, importlib.metadata, itertools, re, functools, os
+from tinygrad.helpers import flatten, unwrap
+from clang.cindex import Config, Index, CursorKind as CK, TranslationUnit as TU, LinkageKind as LK, TokenKind as ToK, TypeKind as TK
+from clang.cindex import PrintingPolicy as PP, PrintingPolicyProperty as PPP, SourceRange
+
+assert importlib.metadata.version('clang')[:2] == "20"
+if not Config.loaded: Config.set_library_file(os.getenv("LIBCLANG_PATH", ctypes.util.find_library("clang-20")))
+
+def fst(c): return next(c.get_children())
+def last(c): return list(c.get_children())[-1]
+def readext(f, fst, snd=None):
+  with open(f, "r") as f:
+    f.seek(start:=(fst.start.offset if isinstance(fst, SourceRange) else fst))
+    return f.read((fst.end.offset if isinstance(fst, SourceRange) else snd)-start)
+def attrs(c): return list(filter(lambda k: (v:=k.value) >= 400 and v < 500, map(lambda c: c.kind, c.get_children())))
+
+base_rules = [(r'\s*\\\n\s*', ' '), (r'\s*\n\s*', ' '), (r'//.*', ''), (r'/\*.*?\*/', ''), (r'\b(0[xX][0-9a-fA-F]+|\d+)[uUlL]+\b', r'\1'),
+              (r'\b0+(?=\d)', ''), (r'\s*&&\s*', r' and '), (r'\s*\|\|\s*', r' or '), (r'\s*!\s*', ' not '),
+              (r'(struct|union|enum)\s*([a-zA-Z_][a-zA-Z0-9_]*\b)', r'\1_\2'),
+              (r'\((unsigned )?(char|uint64_t)\)', ''), (r'^.*\d+:\d+.*$', ''), (r'^.*\w##\w.*$', '')]
+
+ints = (TK.INT, TK.UINT, TK.LONG, TK.ULONG, TK.LONGLONG, TK.ULONGLONG)
+
+def gen(dll, files, args=[], prolog=[], rules=[], epilog=[], recsym=False, use_errno=False, anon_names={}, types={}, parse_macros=True):
+  macros, lines, anoncnt, types = [], [], itertools.count().__next__, {k:(v,True) for k,v in types.items()}
+  def tname(t, suggested_name=None, typedef=None) -> str:
+    suggested_name = anon_names.get(f"{(decl:=t.get_declaration()).location.file}:{decl.location.line}", suggested_name)
+    nonlocal lines, types, anoncnt
+    tmap = {TK.VOID:"None", TK.CHAR_U:"ctypes.c_ubyte", TK.UCHAR:"ctypes.c_ubyte", TK.CHAR_S:"ctypes.c_char", TK.SCHAR:"ctypes.c_char",
+            **{getattr(TK, k):f"ctypes.c_{k.lower()}" for k in ["BOOL", "WCHAR", "FLOAT", "DOUBLE", "LONGDOUBLE"]},
+            **{getattr(TK, k):f"ctypes.c_{'u' if 'U' in k else ''}int{sz}" for sz,k in
+               [(16, "USHORT"), (16, "SHORT"), (32, "UINT"), (32, "INT"), (64, "ULONG"), (64, "LONG"), (64, "ULONGLONG"), (64, "LONGLONG")]}}
+
+    if t.kind in tmap: return tmap[t.kind]
+    if t.spelling in types and types[t.spelling][1]: return types[t.spelling][0]
+    if ((f:=t).kind in (fks:=(TK.FUNCTIONPROTO, TK.FUNCTIONNOPROTO))) or (t.kind == TK.POINTER and (f:=t.get_pointee()).kind in fks):
+      return f"ctypes.CFUNCTYPE({tname(f.get_result())}{(', '+', '.join(map(tname, f.argument_types()))) if f.kind==TK.FUNCTIONPROTO else ''})"
+    match t.kind:
+      case TK.POINTER: return "ctypes.c_void_p" if (ptr:=t.get_pointee()).kind == TK.VOID else f"ctypes.POINTER({tname(ptr)})"
+      case TK.ELABORATED: return tname(t.get_named_type(), suggested_name)
+      case TK.TYPEDEF if t.spelling == t.get_canonical().spelling: return tname(t.get_canonical())
+      case TK.TYPEDEF:
+        defined, nm = (canon:=t.get_canonical()).spelling in types, tname(canon, typedef=t.spelling.replace('::', '_'))
+        types[t.spelling] = nm if t.spelling.startswith("__") else t.spelling.replace('::', '_'), True
+        # RECORDs need to handle typedefs specially to allow for self-reference
+        if canon.kind != TK.RECORD or defined: lines.append(f"{t.spelling.replace('::', '_')} = {nm}")
+        return types[t.spelling][0]
+      case TK.RECORD:
+        # TODO: packed unions
+        # TODO: pragma pack support
+        # check for forward declaration
+        if t.spelling in types: types[t.spelling] = (nm:=types[t.spelling][0]), len(list(t.get_fields())) != 0
+        else:
+          if decl.is_anonymous():
+            types[t.spelling] = (nm:=(suggested_name or (f"_anon{'struct' if decl.kind == CK.STRUCT_DECL else 'union'}{anoncnt()}")), True)
+          else: types[t.spelling] = (nm:=t.spelling.replace(' ', '_').replace('::', '_')), len(list(t.get_fields())) != 0
+          lines.append(f"class {nm}({'Struct' if decl.kind==CK.STRUCT_DECL else 'ctypes.Union'}): pass")
+          if typedef: lines.append(f"{typedef} = {nm}")
+        acnt = itertools.count().__next__
+        ll=["  ("+((fn:=f"'_{acnt()}'")+f", {tname(f.type, nm+fn[1:-1])}" if f.is_anonymous_record_decl() else f"'{f.spelling}', "+
+            tname(f.type, f'{nm}_{f.spelling}'))+(f',{f.get_bitfield_width()}' if f.is_bitfield() else '')+")," for f in t.get_fields()]
+        lines.extend(([f"{nm}._anonymous_ = ["+", ".join(f"'_{i}'" for i in range(n))+"]"] if (n:=acnt()) else [])+
+                     ([f"{nm}._packed_ = True"] * (CK.PACKED_ATTR in attrs(decl)))+([f"{nm}._fields_ = [",*ll,"]"] if ll else []))
+        return nm
+      case TK.ENUM:
+        # TODO: C++ and GNU C have forward declared enums
+        if decl.is_anonymous(): types[t.spelling] = suggested_name or f"_anonenum{anoncnt()}", True
+        else: types[t.spelling] = t.spelling.replace(' ', '_').replace('::', '_'), True
+        lines.append(f"{types[t.spelling][0]} = CEnum({tname(decl.enum_type)})\n" +
+                     "\n".join(f"{e.spelling} = {types[t.spelling][0]}.define('{e.spelling}', {e.enum_value})" for e in decl.get_children()
+                     if e.kind == CK.ENUM_CONSTANT_DECL) + "\n")
+        return types[t.spelling][0]
+      case TK.CONSTANTARRAY:
+        return f"({tname(t.get_array_element_type(), suggested_name.rstrip('s') if suggested_name else None)} * {t.get_array_size()})"
+      case TK.INCOMPLETEARRAY: return f"({tname(t.get_array_element_type(), suggested_name.rstrip('s') if suggested_name else None)} * 0)"
+      case _: raise NotImplementedError(f"unsupported type {t.kind}")
+
+  for f in files:
+    tu = Index.create().parse(f, args, options=TU.PARSE_DETAILED_PROCESSING_RECORD)
+    (pp:=PP.create(tu.cursor)).set_property(PPP.TerseOutput, 1)
+    for c in tu.cursor.walk_preorder():
+      if str(c.location.file) != str(f) and (not recsym or c.kind not in (CK.FUNCTION_DECL,)): continue
+      rollback = lines, types
+      try:
+        match c.kind:
+          case CK.FUNCTION_DECL if c.linkage == LK.EXTERNAL and dll:
+            # TODO: we could support name-mangling
+            lines.append(f"# {c.pretty_printed(pp)}\ntry: ({c.spelling}:=dll.{c.spelling}).restype, {c.spelling}.argtypes = "
+              f"{tname(c.result_type)}, [{', '.join(tname(arg.type) for arg in c.get_arguments())}]\nexcept AttributeError: pass\n")
+          case CK.STRUCT_DECL | CK.UNION_DECL | CK.TYPEDEF_DECL | CK.ENUM_DECL: tname(c.type)
+          case CK.MACRO_DEFINITION if parse_macros and len(toks:=list(c.get_tokens())) > 1:
+            if toks[1].spelling == '(' and toks[0].extent.end.column == toks[1].extent.start.column:
+              it = iter(toks[1:])
+              _args = [t.spelling for t in itertools.takewhile(lambda t:t.spelling!=')', it) if t.kind == ToK.IDENTIFIER]
+              if len(body:=list(it)) == 0: continue
+              macros += [f"{c.spelling} = lambda {','.join(_args)}: {readext(f, body[0].location.offset, toks[-1].extent.end.offset)}"]
+            else: macros += [f"{c.spelling} = {readext(f, toks[1].location.offset, toks[-1].extent.end.offset)}"]
+          case CK.VAR_DECL if c.linkage == LK.INTERNAL:
+            if (c.type.kind == TK.CONSTANTARRAY and c.type.get_array_element_type().get_canonical().kind in ints and
+                (init:=last(c)).kind == CK.INIT_LIST_EXPR and all(re.match(r"\[.*\].*=", readext(f, c.extent)) for c in init.get_children())):
+              cs = init.get_children()
+              macros += [f"{c.spelling} = {{{','.join(f'{readext(f,next(it:=c.get_children()).extent)}:{readext(f,next(it).extent)}' for c in cs)}}}"]
+            elif c.type.get_canonical().kind in ints: macros += [f"{c.spelling} = {readext(f, last(c).extent)}"]
+            else: macros += [f"{c.spelling} = {tname(c.type)}({readext(f, last(c).extent)})"]
+          case CK.VAR_DECL if c.linkage == LK.EXTERNAL and dll:
+            lines.append(f"try: {c.spelling} = {tname(c.type)}.in_dll(dll, '{c.spelling}')\nexcept (ValueError,AttributeError): pass")
+      except NotImplementedError as e:
+        print(f"skipping {c.spelling}: {e}")
+        lines, types = rollback
+  main = (f"# mypy: ignore-errors\nimport ctypes{', os' if any('os' in s for s in dll) else ''}\n"
+    "from tinygrad.helpers import unwrap\nfrom tinygrad.runtime.support.c import Struct, CEnum, _IO, _IOW, _IOR, _IOWR\n" + '\n'.join([*prolog,
+      *(["from ctypes.util import find_library"]*any('find_library' in s for s in dll)),
+      *(["def dll():",*flatten([[f"  try: return ctypes.CDLL(unwrap({d}){', use_errno=True' if use_errno else ''})",'  except: pass'] for d in dll]),
+         "  return None", "dll = dll()\n"]*bool(dll)), *lines]) + '\n')
+  macros = [r for m in macros if (r:=functools.reduce(lambda s,r:re.sub(r[0], r[1], s), rules + base_rules, m))]
+  while True:
+    try:
+      exec(main + '\n'.join(macros), {})
+      break
+    except (SyntaxError, NameError, TypeError) as e:
+      macrono = unwrap(e.lineno if isinstance(e, SyntaxError) else unwrap(unwrap(e.__traceback__).tb_next).tb_lineno) - main.count('\n') - 1
+      assert macrono >= 0 and macrono < len(macros), f"error outside macro range: {e}"
+      print(f"skipping {macros[macrono]}: {e}")
+      del macros[macrono]
+    except Exception as e: raise Exception("parsing failed") from e
+  return main + '\n'.join(macros + epilog)
@@ -0,0 +1,72 @@
+import ctypes, functools, sys
+from typing import TYPE_CHECKING
+
+def _do_ioctl(__idir, __base, __nr, __struct, __fd, **kwargs):
+  import tinygrad.runtime.support.hcq as hcq, fcntl
+  ioctl = __fd.ioctl if isinstance(__fd, hcq.FileIOInterface) else functools.partial(fcntl.ioctl, __fd)
+  if (rc:=ioctl((__idir<<30)|(ctypes.sizeof(out:=__struct(**kwargs))<<16)|(__base<<8)|__nr, out)): raise RuntimeError(f"ioctl returned {rc}")
+  return out
+
+def _IO(base, nr): return functools.partial(_do_ioctl, 0, ord(base) if isinstance(base, str) else base, nr, None)
+def _IOW(base, nr, typ): return functools.partial(_do_ioctl, 1, ord(base) if isinstance(base, str) else base, nr, typ)
+def _IOR(base, nr, typ): return functools.partial(_do_ioctl, 2, ord(base) if isinstance(base, str) else base, nr, typ)
+def _IOWR(base, nr, typ): return functools.partial(_do_ioctl, 3, ord(base) if isinstance(base, str) else base, nr, typ)
+
+def CEnum(typ: type[ctypes._SimpleCData]):
+  class _CEnum(typ): # type: ignore
+    _val_to_name_: dict[int,str] = {}
+
+    @classmethod
+    def from_param(cls, val): return val if isinstance(val, cls) else cls(val)
+    @classmethod
+    def get(cls, val, default="unknown"): return cls._val_to_name_.get(val.value if isinstance(val, cls) else val, default)
+    @classmethod
+    def items(cls): return cls._val_to_name_.items()
+    @classmethod
+    def define(cls, name, val):
+      cls._val_to_name_[val] = name
+      return val
+
+    def __eq__(self, other): return self.value == other
+    def __repr__(self): return self.get(self) if self.value in self.__class__._val_to_name_ else str(self.value)
+
+  return _CEnum
+
+# supports gcc (C11) __attribute__((packed))
+if TYPE_CHECKING: Struct = ctypes.Structure
+else:
+  class MetaStruct(type(ctypes.Structure)):
+    def __new__(mcs, name, bases, dct):
+      fields = dct.pop("_fields_", None)
+      cls = super().__new__(mcs, name, bases, dct)
+      if dct.get("_packed_", False) and fields is not None: mcs._build(cls, fields)
+      return cls
+
+    def __setattr__(cls, k, v):
+      # NB: _fields_ must be set after _packed_ because PyCStructType_setattro marks _fields_ as final.
+      if k == "_fields_" and getattr(cls, "_packed_", False): type(cls)._build(cls, v)
+      elif k == "_packed_" and hasattr(cls, "_fields_"): type(cls)._build(cls, cls._fields_)
+      else: super().__setattr__(k, v)
+
+    @staticmethod
+    def _build(cls, fields):
+      o = 0
+      for n,t,b in [(f[0], f[1], f[2] if len(f) == 3 else 0) for f in fields]:
+        if b == 0: o = (o + 7) & ~7
+        m = (1 << (sz:=ctypes.sizeof(t)*8 if b == 0 else b)) - 1
+        def _s(self,v,m,s,b): self._data[:] = ((int.from_bytes(self._data,sys.byteorder)&~(m<<s))|((v&m)<<s)).to_bytes(len(self._data), sys.byteorder)
+        setattr(cls, n, property(functools.partial(lambda self,m,s:(int.from_bytes(self._data,sys.byteorder)>>s)&m,m=m,s=o),
+                                 functools.partial(_s,m=m,s=o,b=b)))
+        o += sz
+
+      type(ctypes.Structure).__setattr__(cls, '_fields_', [('_data', ctypes.c_ubyte * ((o + 7) // 8))])
+      type(ctypes.Structure).__setattr__(cls, '_packed_', True)
+      setattr(cls, '_packed_fields_', fields)
+
+  class Struct(ctypes.Structure, metaclass=MetaStruct):
+    def __init__(self, *args, **kwargs):
+      if hasattr(self, '_packed_fields_'):
+        for f,v in zip(self._packed_fields_, args): setattr(self, f[0], v)
+        for k,v in kwargs.items(): setattr(self, k, v)
+      else: super().__init__(*args, **kwargs)
+
@@ -1,4 +1,5 @@
-import ctypes, subprocess
+import ctypes
+from tinygrad.helpers import system
 import tinygrad.runtime.autogen.comgr as comgr
 assert comgr.AMD_COMGR_LANGUAGE_HIP == 4
 try:
@@ -13,7 +14,7 @@ from tinygrad.runtime.support.compiler_cpu import LLVMCompiler
 from tinygrad.helpers import OSX, to_char_p_p

 def amdgpu_disassemble(lib:bytes):
-  asm = subprocess.check_output(["llvm-objdump" if OSX else "/opt/rocm/llvm/bin/llvm-objdump", '-d', '-'], input=lib).decode("utf-8").splitlines()
+  asm = system(f"{'llvm-objdump' if OSX else '/opt/rocm/llvm/bin/llvm-objdump'} -d -", input=lib).splitlines()
  while asm and ("s_nop 0" in asm[-1] or "s_code_end" in asm[-1]): asm.pop()
  print("\n".join(asm))

@@ -1,6 +1,6 @@
 import subprocess, hashlib, tempfile, ctypes, re, pathlib
 from typing import Callable
-from tinygrad.helpers import to_char_p_p, colored, init_c_var, getenv
+from tinygrad.helpers import to_char_p_p, colored, init_c_var, getenv, system
 import tinygrad.runtime.autogen.nvrtc as nvrtc
 from tinygrad.device import Compiler, CompileError

@@ -37,7 +37,7 @@ def cuda_disassemble(lib:bytes, arch:str):
    fn = (pathlib.Path(tempfile.gettempdir()) / f"tinycuda_{hashlib.md5(lib).hexdigest()}").as_posix()
    with open(fn, "wb") as f: f.write(lib)
    subprocess.run(["ptxas", f"-arch={arch}", "-o", fn, fn], check=False, stderr=subprocess.DEVNULL) # optional ptx -> sass step for CUDA=1
-    print(subprocess.check_output(['nvdisasm', fn]).decode('utf-8'))
+    print(system(f'nvdisasm {fn}'))
  except Exception as e: print("Failed to generate SASS", str(e), "Make sure your PATH contains ptxas/nvdisasm binary of compatible version.")

 class CUDACompiler(Compiler):
@@ -1,6 +1,6 @@
-import base64, ctypes, pathlib, tempfile, hashlib, subprocess
+import base64, ctypes, pathlib, tempfile, hashlib
 from tinygrad.device import Compiler
-from tinygrad.helpers import cpu_objdump
+from tinygrad.helpers import cpu_objdump, system
 import tinygrad.runtime.autogen.mesa as mesa
 from tinygrad.runtime.support.compiler_cpu import CPULLVMCompiler, expect, cerr
 try: import tinygrad.runtime.autogen.llvm as llvm
@@ -82,5 +82,5 @@ class NAKCompiler(NIRCompiler):
    try:
      fn = (pathlib.Path(tempfile.gettempdir()) / f"tinynak_{hashlib.md5(lib).hexdigest()}").as_posix()
      with open(fn, "wb") as f: f.write(lib[ctypes.sizeof(mesa.struct_nak_shader_info):])
-      print(subprocess.check_output(['nvdisasm', "-b", f"SM{self.arch[3:]}", fn]).decode('utf-8'))
+      print(system(f"nvdisasm -b SM{self.arch[3:]} {fn}"))
    except Exception as e: print("Failed to generate SASS", str(e), "Make sure your PATH contains nvdisasm binary of compatible version.")
@@ -1,7 +1,7 @@
 import struct, ctypes, ctypes.util
 from dataclasses import dataclass
 from tinygrad.helpers import getbits, i2u, unwrap
-import tinygrad.runtime.autogen.libc as libc
+from tinygrad.runtime.autogen import libc

@dataclass(frozen=True)
 class ElfSection: name:str; header:libc.Elf64_Shdr; content:bytes # noqa: E702
@@ -1,5 +1,5 @@
-import ctypes.util, os, sys, subprocess
-from tinygrad.helpers import DEBUG, OSX, getenv
+import ctypes.util, os, sys
+from tinygrad.helpers import DEBUG, OSX, getenv, system

 if sys.platform == 'win32':
  # Windows llvm distribution doesn't seem to add itself to PATH or anywhere else where it can be easily retrieved from.
@@ -10,7 +10,7 @@ if sys.platform == 'win32':
 elif OSX:
  # Will raise FileNotFoundError if brew is not installed
  # `brew --prefix` will return even if formula is not installed
-  if not os.path.exists(brew_prefix:=subprocess.check_output(['brew', '--prefix', 'llvm@20']).decode().strip()):
+  if not os.path.exists(brew_prefix:=system("brew --prefix llvm@20")):
    raise FileNotFoundError('LLVM not found, you can install it with `brew install llvm@20`')
  LLVM_PATH: str|None = os.path.join(brew_prefix, 'lib', 'libLLVM.dylib')
 else:
@@ -49,8 +49,9 @@ class NVRpcQueue:
    self.seq += 1
    self.gsp.nvdev.NV_PGSP_QUEUE_HEAD[0].write(0x0)

-  def wait_resp(self, cmd:int) -> memoryview:
-    while True:
+  def wait_resp(self, cmd:int, timeout=10000) -> memoryview:
+    start_time = int(time.perf_counter() * 1000)
+    while (int(time.perf_counter() * 1000) - start_time) < timeout:
      System.memory_barrier()
      if self.rx.readPtr == self.tx.writePtr: continue

@@ -73,6 +74,7 @@ class NVRpcQueue:

      if hdr.rpc_result != 0: raise RuntimeError(f"RPC call {hdr.function} failed with result {hdr.rpc_result}")
      if hdr.function == cmd: return msg
+    raise RuntimeError(f"Timeout waiting for RPC response for command {cmd}")

 class NV_FLCN(NV_IP):
  def init_sw(self):
@@ -12,6 +12,11 @@ MAP_FIXED, MAP_LOCKED, MAP_POPULATE, MAP_NORESERVE = 0x10, 0 if OSX else 0x2000,
 class PCIBarInfo: addr:int; size:int # noqa: E702

 class _System:
+  def write_sysfs(self, path:str, value:str, msg:str, expected:str|None=None):
+    if FileIOInterface(path, os.O_RDONLY).read().splitlines()[0] != (expected or value):
+      os.system(cmd:=f"sudo sh -c 'echo {value} > {path}'")
+      if FileIOInterface(path, os.O_RDONLY).read().splitlines()[0] != (expected or value): raise RuntimeError(f"{msg}. Please run {cmd} manually.")
+
  @functools.cached_property
  def atomic_lib(self): return ctypes.CDLL(ctypes.util.find_library('atomic')) if sys.platform == "linux" else None

@@ -26,9 +31,7 @@ class _System:

  @functools.cached_property
  def pagemap(self) -> FileIOInterface:
-    if FileIOInterface(reloc_sysfs:="/proc/sys/vm/compact_unevictable_allowed", os.O_RDONLY).read()[0] != "0":
-      os.system(cmd:=f"sudo sh -c 'echo 0 > {reloc_sysfs}'")
-      assert FileIOInterface(reloc_sysfs, os.O_RDONLY).read()[0] == "0", f"Failed to disable migration of locked pages. Please run {cmd} manually."
+    self.write_sysfs("/proc/sys/vm/compact_unevictable_allowed", "0", "Failed to disable migration of locked pages")
    return FileIOInterface("/proc/self/pagemap", os.O_RDONLY)

  @functools.cached_property
@@ -90,12 +93,12 @@ class _System:
    if data is not None: sysmem_view[:len(data)] = data
    return sysmem_view, [p + i for p, sz in paddrs for i in range(0, sz, 0x1000)][:ceildiv(size, 0x1000)]

-  def pci_scan_bus(self, target_vendor:int, target_devices:list[int]) -> list[str]:
+  def pci_scan_bus(self, target_vendor:int, target_devices:list[tuple[int, list[int]]]) -> list[str]:
    result = []
    for pcibus in FileIOInterface("/sys/bus/pci/devices").listdir():
      vendor = int(FileIOInterface(f"/sys/bus/pci/devices/{pcibus}/vendor").read(), 16)
      device = int(FileIOInterface(f"/sys/bus/pci/devices/{pcibus}/device").read(), 16)
-      if vendor == target_vendor and device in target_devices: result.append(pcibus)
+      if vendor == target_vendor and any((device & mask) in devlist for mask, devlist in target_devices): result.append(pcibus)
    return sorted(result)

  def pci_setup_usb_bars(self, usb:ASM24Controller, gpu_bus:int, mem_base:int, pref_mem_base:int) -> dict[int, PCIBarInfo]:
@@ -244,7 +247,7 @@ class LNXPCIIfaceBase:
  dev_impl:PCIDevImplBase
  gpus:ClassVar[list[str]] = []

-  def __init__(self, dev, dev_id, vendor, devices, bars, vram_bar, va_start, va_size):
+  def __init__(self, dev, dev_id, vendor, devices:list[tuple[int, list[int]]], bars, vram_bar, va_start, va_size):
    if len((cls:=type(self)).gpus) == 0:
      cls.gpus = hcq_filter_visible_devices(System.pci_scan_bus(vendor, devices))

@@ -1,10 +1,10 @@
-import ctypes.util, os, subprocess, platform, sysconfig
-from tinygrad.helpers import OSX
+import ctypes.util, os, platform, sysconfig
+from tinygrad.helpers import system, OSX

 WEBGPU_PATH: str | None

 if OSX:
-  if not os.path.exists(brew_prefix:=subprocess.check_output(['brew', '--prefix', 'dawn']).decode().strip()):
+  if not os.path.exists(brew_prefix:=system("brew --prefix dawn")):
    raise FileNotFoundError('dawn library not found. Install it with `brew tap wpmed92/dawn && brew install dawn`')
  WEBGPU_PATH = os.path.join(brew_prefix, 'lib', 'libwebgpu_dawn.dylib')
 elif platform.system() == "Windows":
@@ -16,9 +16,6 @@ def realize_srcs(ctx:dict[UOp, None], rb:UOp) -> None:
  for s in rb.src:
    if s.base.op not in ALWAYS_CONTIGUOUS: ctx[s] = None

-def realize_store(ctx:dict[UOp, None], a:UOp) -> None:
-  if a.src[1].op not in ALWAYS_CONTIGUOUS: ctx[a.src[1]] = None
-
 def realize_assign(ctx:dict[UOp, None], a:UOp) -> None:
  if a.src[1].op not in ALWAYS_CONTIGUOUS: ctx[a.src[1]] = None
  # if it's a kernel, we don't realize it
@@ -33,8 +30,6 @@ pm_generate_realize_map = PatternMatcher([
  (UPat((Ops.COPY, Ops.MSELECT, Ops.MSTACK), name="rb"), realize_srcs),
  # realize ASSIGN and input to assign (might be optimized out)
  (UPat(Ops.ASSIGN, name="a"), realize_assign),
-  # realize STORE
-  (UPat(Ops.STORE, name="a"), realize_store),
 ])

@dataclass(frozen=True)
@@ -55,14 +50,13 @@ class IndexingContext:
    # if a range has a 1 src, it's the same as UOp.const(dtypes.index, 0)
    return UOp.range(s, next(self.range_idx), axistype) if resolve(s!=1) else UOp.const(dtypes.index, 0)

-ops_allowed_after = (Ops.KERNEL, Ops.RANGE)
 def create_bufferize_and_index_based_on_ranges(ctx:IndexingContext, x:UOp):
  if x.op in {Ops.BUFFERIZE, Ops.INDEX}: return None
-  if x.op is Ops.AFTER and x.src[1].op in ops_allowed_after: return None
+  if x.op is Ops.AFTER and x.src[1].op is Ops.KERNEL: return None
  new_srcs = []
  for s in x.src:
    new_src = s
-    if s.op in {Ops.BUFFER, Ops.BUFFER_VIEW, Ops.MSTACK, Ops.MSELECT} or (s.op is Ops.AFTER and s.src[1].op in ops_allowed_after):
+    if s.op in {Ops.BUFFER, Ops.BUFFER_VIEW, Ops.MSTACK, Ops.MSELECT} or (s.op is Ops.AFTER and s.src[1].op is Ops.KERNEL):
      if x in ctx.range_map: new_src = new_src.index(*ctx.range_map[x][0])
    elif s in ctx.realize_map:
      realized_ranges = ctx.realize_map[s]
@@ -182,7 +176,7 @@ def run_rangeify(tsink:UOp, debug:bool=False) -> tuple[UOp, IndexingContext]:
      # mark all ranges as ended
      assert rctx.realize_map[x] is None
      rctx.realize_map[x] = list(range(len(x.shape)))
-    elif x.op in {Ops.MSTACK, Ops.MSELECT, Ops.END}:
+    elif x.op in {Ops.MSTACK, Ops.MSELECT}:
      # treat MSTACK/MSELECT like SINK
      continue
    elif len(consumer_rngs) == 0:
@@ -258,7 +252,7 @@ def run_rangeify(tsink:UOp, debug:bool=False) -> tuple[UOp, IndexingContext]:
      else:
        disp = render_ranges(rngs, out_rngs, realized=realized_ranges)
      print("***" if x in rctx.realize_map else "   ",
-            f"{len(consumer_map[x]):2d} {str(x.op):20s} {str(x.shape):35s} {len(ending_ranges[x]):2d}", disp)
+            f"{len(consumer_map[x]):2d} {str(x.op):20s} {str(x._shape):35s} {len(ending_ranges[x]):2d}", disp)

    # assign to the range map. rngs are the input ranges, out_rngs are the output ranges, from the x op.
    rctx.range_map[x] = (rngs, out_rngs)
@@ -5,7 +5,7 @@ from tinygrad.uop.ops import PatternMatcher, UPat, Ops, UOp, resolve, GroupOp, _
 from tinygrad.uop.ops import track_rewrites, graph_rewrite, identity_element, sint, AxisType, BottomUpGate, Kernel, _remove_all_tags
 from tinygrad.uop.symbolic import symbolic
 from tinygrad.helpers import argsort, prod, all_same, pluralize, getenv, flatten, dedup, all_int, DEBUG, SPLIT_REDUCEOP, DEBUG_RANGEIFY
-from tinygrad.helpers import PCONTIG, partition, get_single_element, unwrap
+from tinygrad.helpers import PCONTIG, partition, get_single_element, unwrap, disable_gc
 from tinygrad.codegen.simplify import pm_flatten_range, pm_reduce_simplify
 from tinygrad.codegen.opt import Opt
 from tinygrad.schedule.indexing import run_rangeify, BufferizeOpts, ALWAYS_CONTIGUOUS, IndexingContext, apply_movement_op
@@ -396,7 +396,6 @@ def handle_after(ctx:LocalAddBufferContext, after:UOp):
  return buf

 def renumber_range(ctx:LocalAddBufferContext, r:UOp):
-  if r.arg[-1] == AxisType.OUTER: return None
  if r.tag != (): return None
  ret = r.replace(arg=(ctx.range,)+r.arg[1:], tag=None)
  ctx.range += 1
@@ -470,10 +469,7 @@ pm_add_range_tags = PatternMatcher([
 ])

 def split_store(ctx:list[UOp], x:UOp) -> UOp|None:
-  if len([r for r in x.ranges if r.arg[-1] != AxisType.OUTER]): return None
-
-  # ends of outer range don't go in kernels
-  if x.op is Ops.END and x.src[1].op is Ops.RANGE and x.src[1].arg[-1] == AxisType.OUTER: return None
+  if len(x.ranges): return None

  # local kernel rewrite
  lctx = LocalAddBufferContext()
@@ -529,6 +525,7 @@ replace_contiguous = PatternMatcher([
  (UPat(GroupOp.ALU, name="alu"), lambda ctx,alu: alu.replace(src=new_src) if (new_src:=tuple(ctx.get(s, s) for s in alu.src)) != alu.src else None),
 ])

+@disable_gc()
@track_rewrites(lambda _,ret: f"Schedule {pluralize('Kernel', len([u for u in UOp.sink(*ret.values()).toposort() if u.op is Ops.KERNEL]))}", True)
 def get_rangeify_map(sink:UOp) -> dict[UOp, UOp]:
  if getenv("VIZ"): graph_rewrite(sink, PatternMatcher([]), name="View Input Graph")
@@ -172,7 +172,7 @@ class Tensor(OpMixin):

  def _apply_uop(self, fxn:Callable, *x:Tensor, extra_args=(), **kwargs) -> Tensor:
    new_uop: UOp = fxn(*[t.uop for t in (self,)+x], *extra_args, **kwargs)
-    if (metadata:=_METADATA.get()) is not None: all_metadata[new_uop] = (metadata,)
+    if (metadata:=_METADATA.get()) is not None and TRACEMETA >= 1: all_metadata[new_uop] = (metadata,)
    needs_input_grad = [t.requires_grad for t in (self,)+x]
    return Tensor(new_uop, device=new_uop.device, requires_grad=True if any(needs_input_grad) else None if None in needs_input_grad else False)

@@ -2093,25 +2093,6 @@ class Tensor(OpMixin):

  # ***** processing ops *****

-  def _pool(self, k_:tuple[sint, ...], stride:int|tuple[int, ...]=1, dilation:int|tuple[int, ...]=1) -> Tensor:
-    assert len(self.shape) >= len(k_), f"can't pool {self.shape} with {k_}"
-    s_, d_ = make_tuple(stride, len(k_)), make_tuple(dilation, len(k_))
-    assert len(k_) == len(s_) == len(d_), f"stride/dilation mismatch kernel:{k_} stride:{s_} dilation:{d_}"
-    noop, i_ = [None] * (self.ndim-len(k_)), self.shape[-len(k_):]
-    assert all(resolve(d*(k-1)+1 <= i) for k,d,i in zip(k_,d_,i_)), "kernel size cannot be greater than actual input size"
-    o_ = [ceildiv(i-d*(k-1), s) for i,d,k,s in zip(i_,d_,k_,s_)]
-    # input size scaling factor to make sure shrink for stride is possible
-    f_ = [smax(1, ceildiv(o*s - d, i)) for o,s,i,d in zip(o_,s_,i_,d_)]
-    # repeats such that we don't need padding
-    x = self.repeat([1]*len(noop) + [ceildiv(k*(i*f+d),i) for k,i,d,f in zip(k_,i_,d_,f_)])
-    # handle dilation
-    x = x.shrink_to(noop + [k*(i*f+d) for k,i,d,f in zip(k_,i_,d_,f_)]).reshape(noop + flatten((k,(i*f+d)) for k,i,d,f in zip(k_,i_,d_,f_)))
-    # handle stride
-    x = x.shrink_to(noop + flatten((k,o*s) for k,o,s in zip(k_,o_,s_))).reshape(noop + flatten((k,o,s) for k,o,s in zip(k_,o_,s_)))
-    x = x.shrink_to(noop + flatten((k,o,1) for k,o in zip(k_,o_))).reshape(noop + flatten((k,o) for k,o in zip(k_,o_)))
-    # permute to move reduce to the end
-    return x.permute(*range(len(noop)), *[len(noop)+i*2+1 for i in range(len(i_))], *[len(noop)+i*2 for i in range(len(i_))])
-
  def _resolve_pool_pads(self, padding:int|Sequence[int], dims:int) -> Sequence[int]:
    if not isinstance(padding, int) and not (len(padding) == 2*dims or len(padding) == dims):
      raise ValueError(f"Padding must be an int or a sequence of length {dims} or {2*dims}, but got {padding=} for {self.shape=} with {dims=}.")
@@ -4111,18 +4092,17 @@ class Tensor(OpMixin):

  def image_dot(self, w:Tensor, dtype:DTypeLike|None=None) -> Tensor:
    # NOTE: we use a 1x1 conv2d to do the matmul. mxk @ kxn = (1,k,m,1).conv2d(n,k,1,1)
-    x, dx, dw = self, self.ndim, w.ndim
-    if not (dx > 0 and dw > 0): raise RuntimeError(f"both tensors need to be at least 1D, got {dx}D and {dw}D")
-    if x.shape[-1] != w.shape[-min(w.ndim, 2)]: raise RuntimeError(f"cannot image_dot {x.shape} and {w.shape}")
+    if not (self.ndim > 0 and w.ndim > 0): raise RuntimeError(f"both tensors need to be at least 1D, got {self.ndim=}, {w.ndim=}")
+    if self.shape[-1] != w.shape[-min(w.ndim, 2)]: raise RuntimeError(f"cannot image_dot {self.shape} and {w.shape}")

    bs, groups, cin, cout = prod(self.shape[0:-2]), prod(w.shape[0:-2]), w.shape[-2], w.shape[-1]
-    out_shape_t = self.shape[0:-2] + (cout,-1) if len(self.shape) > 1 else (cout, )
+    out_shape_t = self.shape[0:-2] + (cout,-1) if len(self.shape) > 1 else (cout,)

    # NOTE: with NHWC we can remove the transposes
    # bs x groups*cin x H x W
-    cx = self.transpose(self.ndim-1, self.ndim-2).reshape((bs//groups, groups*cin, -1, 1))
+    cx = self.transpose(self.ndim-1, self.ndim-2).reshape(bs//groups, groups*cin, -1, 1)
    # groups*cout x cin x H, W
-    cw = w.transpose(w.ndim-1, w.ndim-2).reshape((groups*cout, cin, 1, 1))
+    cw = w.transpose(w.ndim-1, w.ndim-2).reshape(groups*cout, cin, 1, 1)
    return cx.image_conv2d(cw, groups=groups, dtype=dtype).reshape(out_shape_t).transpose(self.ndim-1, self.ndim-2)

  def image_conv2d(self, weight:Tensor, bias:Tensor|None=None, groups=1, stride=1, dilation=1, padding=0, dtype=None) -> Tensor:
@@ -4135,10 +4115,9 @@ class Tensor(OpMixin):
    if cin % 4 != 0 and not (cin == 1 and groups%4 == 0):
      x = x.reshape(bs, groups, cin, iy, ix)   # do this always?
      added_input_channels = 4 - (cin % 4)
-      w = w.pad(tuple((0, added_input_channels) if i == 2 else None for i in range(w.ndim)))
-      x = x.pad(tuple((0, added_input_channels) if i == 2 else None for i in range(x.ndim)))
      cin = cin + added_input_channels
-      x = x.reshape(bs, groups*cin, iy, ix)
+      w = w.pad_to(None, None, cin, None, None)
+      x = x.pad_to(None, None, cin, None, None).reshape(bs, groups*cin, iy, ix)

    # hack for non multiples of 4 on rcout
    added_output_channels = 0
@@ -4146,7 +4125,7 @@ class Tensor(OpMixin):
      added_output_channels = 4 - (rcout % 4)
      rcout += added_output_channels
      cout = groups * rcout
-      w = w.pad(tuple((0, added_output_channels) if i == 1 else None for i in range(w.ndim)))
+      w = w.pad_to(None, rcout, None, None, None)

    # packed (note: flipping bs and iy would make the auto-padding work)
    x = x.permute(0,2,3,1)
@@ -4199,7 +4178,7 @@ _METADATA: _ContextVar[Metadata|None] = _ContextVar(default=None)

 def _metadata_wrapper(fn: Callable[P, T]) -> Callable[P, T]:
  def _wrapper(*args: P.args, **kwargs: P.kwargs) -> T:
-    if _METADATA.get() is not None: return fn(*args, **kwargs)
+    if TRACEMETA < 1 or _METADATA.get() is not None: return fn(*args, **kwargs)

    if TRACEMETA >= 2:
      caller_frame = sys._getframe(frame := 1)
@@ -4,26 +4,21 @@ import sys, time, functools, itertools, math, operator, hashlib, os, types, pick
 from dataclasses import dataclass
 from enum import Enum, auto
 from tinygrad.uop import Ops, GroupOp
-from tinygrad.mixin import OpMixin
 from tinygrad.dtype import ConstType, ImageDType, dtypes, DType, truncate, PtrDType, least_upper_dtype, Invalid, InvalidType, AddrSpace
 from tinygrad.helpers import ContextVar, all_int, prod, getenv, all_same, Context, partition, temp, unwrap, T, argfix, Metadata, flatten, TRACEMETA
 from tinygrad.helpers import PICKLE_BUFFERS, PROFILE, dedup, cdiv, cmod, diskcache_put, to_function_name, cpu_profile, TracingKey, VIZ, SPEC, CI
-from tinygrad.helpers import strip_parens, colored, ansilen
+from tinygrad.helpers import strip_parens, colored, ansilen, printable
 if TYPE_CHECKING:
  from tinygrad.device import Buffer, MultiBuffer

 class AxisType(Enum):
  def __repr__(self): return str(self)
  GLOBAL = auto(); WARP = auto(); LOCAL = auto(); LOOP = auto(); GROUP_REDUCE = auto(); REDUCE = auto(); UPCAST = auto(); UNROLL = auto() # noqa: E702
-  THREAD = auto(); OUTER = auto() # noqa: E702
+  THREAD = auto()
 axis_letters = {AxisType.GLOBAL: "g", AxisType.THREAD: "t", AxisType.LOCAL: "l", AxisType.WARP: "w", AxisType.LOOP: "L", AxisType.UPCAST: "u",
-                AxisType.GROUP_REDUCE: "G", AxisType.REDUCE: "R", AxisType.UNROLL: "r", AxisType.OUTER: ("O")}
+                AxisType.GROUP_REDUCE: "G", AxisType.REDUCE: "R", AxisType.UNROLL: "r"}
 axis_colors = {AxisType.GLOBAL: "blue", AxisType.THREAD: "BLUE", AxisType.LOCAL: "cyan", AxisType.WARP: "CYAN", AxisType.LOOP: "WHITE",
-               AxisType.UPCAST: "yellow", AxisType.GROUP_REDUCE: "RED", AxisType.REDUCE: "red", AxisType.UNROLL: "magenta", AxisType.OUTER: "green"}
-
-# NOTE: LOCAL and GROUP_REDUCE have the same priority. the order here matters
-axis_to_pos = {AxisType.OUTER: -2, AxisType.LOOP: -1, AxisType.THREAD: 0, AxisType.GLOBAL: 0, AxisType.WARP: 1, AxisType.LOCAL: 2,
-               AxisType.UPCAST: 3, AxisType.GROUP_REDUCE: 2, AxisType.REDUCE: 4, AxisType.UNROLL: 5}
+               AxisType.UPCAST: "yellow", AxisType.GROUP_REDUCE: "RED", AxisType.REDUCE: "red", AxisType.UNROLL: "magenta"}

 range_start = {Ops.BUFFERIZE: 1, Ops.REDUCE: 1, Ops.STORE: 2, Ops.WMMA: 3, Ops.END: 1}

@@ -111,6 +106,9 @@ class recursive_property(property):
        s.__dict__[self.nm] = val = self.fxn(s)
    return val

+# we import this late so we can use resolve/smax in mixins
+from tinygrad.mixin import OpMixin
+
 # NOTE: this should be frozen, but frozen is slower
@dataclass(eq=False, slots=True)
 class UOp(OpMixin, metaclass=UOpMetaClass):
@@ -221,10 +219,6 @@ class UOp(OpMixin, metaclass=UOpMetaClass):

      # ops with custom handling
      case Ops.KERNEL: return self.arg.ast._shape
-      case Ops.STORE:
-        if isinstance(self.dtype, PtrDType): return (self.ptrdtype.size,)
-        if self.dtype is not dtypes.void: return self.src[0].src[0].shape
-        return None

      # TODO: disallow shape changing bitcast
      case Ops.BITCAST:
@@ -276,7 +270,7 @@ class UOp(OpMixin, metaclass=UOpMetaClass):
          return tuple(1 if i in axis_arg else s for i,s in enumerate(ps))

    # elementwise ops keep the shape the same. all inputs with shape must match
-    if self.op in (GroupOp.Elementwise-{Ops.BITCAST}).union({Ops.COPY, Ops.ASSIGN, Ops.NOOP, Ops.GROUP, Ops.SINK, Ops.ALLREDUCE}):
+    if self.op in (GroupOp.Elementwise-{Ops.BITCAST}).union({Ops.COPY, Ops.ASSIGN, Ops.NOOP, Ops.GROUP, Ops.SINK, Ops.ALLREDUCE, Ops.STORE}):
      # TODO: remove this hack for 3 op assign
      input_shapes = [x._shape for x in (self.src[:2] if self.op is Ops.ASSIGN else self.src) if x._shape is not None]
      if len(input_shapes) == 0: return None
@@ -873,14 +867,6 @@ def get_location() -> tuple[str, int]:
    frm = frm.f_back
  return frm.f_code.co_filename, frm.f_lineno

-@functools.cache
-def lines(fn) -> list[str]:
-  with open(fn) as f: return f.readlines()
-
-def printable(loc:tuple[str, int]) -> str:
-  try: return lines(loc[0])[loc[1]-1].strip()
-  except FileNotFoundError: return "<missing>"
-
 class UPat(OpMixin):
  __slots__ = ("op", "dtype", "arg", "name", "src")
  def __init__(self, op:Ops|tuple[Ops, ...]|set[Ops]|None=None, dtype:DType|tuple[DType, ...]|None=None,
@@ -1000,11 +986,18 @@ def upat_interpret(p:UPat, fxn:Callable) -> Callable:
      return None
  return universal_match

+def fixup_pm_function(fxn) -> Callable:
+  if isinstance(fxn, UPat):
+    # TODO: write this
+    raise NotImplementedError("rhs UPat is not supported")
+  if isinstance(fxn, tuple): return types.FunctionType(*fxn)
+  return fxn
+
 class PatternMatcher:
  def __init__(self, patterns:Sequence[tuple[UPat, Callable|tuple]], compiled=bool(getenv("UPAT_COMPILE", 1))):
    if compiled: from tinygrad.uop.upat import upat_compile
    # if this comes from a pickle, we reconstruct the lambda functions here
-    self.patterns:list[tuple[UPat, Callable]] = [(p,types.FunctionType(*fxn) if isinstance(fxn, tuple) else fxn) for p,fxn in patterns]
+    self.patterns:list[tuple[UPat, Callable]] = [(p,fixup_pm_function(fxn)) for p,fxn in patterns]
    # NOTE: use of DefaultDict here is very dangerous! all keys will live for the lifetime of the PatternMatcher!
    self.pdict: dict[Ops, list[tuple[UPat, Callable, set]]] = {}
    # uop is required, arg is optional
@@ -430,10 +430,10 @@ def uop_given_valid(valid:UOp, uop:UOp, try_simplex=True) -> UOp:
        newuops = [uop.substitute({X:newX}) for X,newX in candidate]
        if any(u is uop for u in newuops): continue  # if any branch doesnt appear in uop, skip
        newuops = [u.simplify().substitute({newX:X}).simplify(full_symbolic=False) for (X,newX),u in zip(candidate,newuops)]
-        if uop.op is Ops.VECTORIZE and len(uop.src) == 2:
+        if all_same(newuops): uop = newuops[0]
+        elif uop.op is Ops.VECTORIZE and len(uop.src) == 2:
          if all_same([uops.src[0] for uops in newuops]): uop = uop.replace(src=(newuops[0].src[0], uop.src[1]))
          if all_same([uops.src[1] for uops in newuops]): uop = uop.replace(src=(uop.src[0], newuops[0].src[1]))
-        elif all_same(newuops): uop = newuops[0]

  # try all the valids together (but only the whole expressions)
  if (s_uop:=uop.substitute(sub_dict:=dict(all_candidates))) is not uop:
@@ -493,7 +493,8 @@ pm_move_where_on_load = PatternMatcher([
 pm_simplify_valid = PatternMatcher([
  # simplify valid
  (UPat(Ops.AND, name="valid"), simplify_valid),
-  (UPat.var("c").where(UPat.var("x", dtype=dtypes.index), invalid_pat), lambda c,x,i: c.where(uop_given_valid(c, x, try_simplex=False), i)),
+  # TODO: this regressed openpilot, not having this regressed cifar
+  # (UPat.var("c").where(UPat.var("x", dtype=dtypes.index), invalid_pat), lambda c,x,i: c.where(uop_given_valid(c, x, try_simplex=False), i)),
 ])

 # this is symbolic 2.0
@@ -292,6 +292,7 @@
  }
  .raw-text > pre {
    display: inline-block;
+    min-width: 100%;
  }
  .raw-text code {
    max-height: none !important;
@@ -25,7 +25,7 @@ const colored = n => d3.create("span").call(s => s.selectAll("span").data(typeof
 const rect = (s) => (typeof s === "string" ? document.querySelector(s) : s).getBoundingClientRect();

 let timeout = null;
-const updateProgress = ({ start }) => {
+const updateProgress = ({ start, err }) => {
  clearTimeout(timeout);
  const msg = document.getElementById("progress-message");
  msg.style.display = "none";
@@ -33,6 +33,11 @@ const updateProgress = ({ start }) => {
    msg.innerText = "Rendering new graph...";
    timeout = setTimeout(() => { msg.style.display = "block"; }, 2000);
  }
+  d3.select("#custom").html("");
+  if (err) {
+    displaySelection("#custom");
+    d3.select("#custom").append(() => d3.create("div").classed("raw-text", true).call(s => s.append(() => codeBlock(err, "txt"))).node());
+  }
 }

 function intersectRect(r1, r2) {
@@ -136,6 +141,10 @@ function renderDag(graph, additions, recenter, layoutOpts) {
    }).attr("class", e => e.value.label.type).attr("id", e => `${e.v}-${e.w}`).datum(e => e.value.label.text));
    if (recenter) document.getElementById("zoom-to-fit-btn").click();
  };
+  worker.onerror = (e) => {
+    e.preventDefault();
+    updateProgress({ err:"Error in graph layout:\n"+e.message });
+  }
 }

 // ** profiler graph
@@ -259,7 +268,10 @@ async function renderProfiler() {
        html.append(() => tabulate([["Name", colored(e.name)], ["Duration", formatTime(e.dur)], ["Start Time", formatTime(e.st)]]).node());
        html.append("div").classed("args", true);
        if (e.info != null) html.append("p").style("white-space", "pre-wrap").text(e.info);
-        if (shapeRef != null) html.append("a").text("View codegen rewrite").on("click", () => switchCtx(shapeRef.ctx, shapeRef.step));
+        if (shapeRef != null) {
+          html.append("a").text("View codegen rewrite").on("click", () => switchCtx(shapeRef.ctx, shapeRef.step));
+          html.append("a").text("View program").on("click", () => switchCtx(shapeRef.ctx, ctxs[shapeRef.ctx+1].steps.findIndex(s => s.name==="View Program")));
+        }
        // tiny device events go straight to the rewrite rule
        const key = k.startsWith("TINY") ? null : `${k}-${j}`;
        if (key != null) shapeMetadata.set(key, html.node());
@@ -539,6 +551,7 @@ document.getElementById("zoom-to-fit-btn").addEventListener("click", () => {

 // **** main VIZ interfacae

+const pathLink = (fp, lineno) => d3.create("a").attr("href", "vscode://file/"+fp+":"+lineno).text(`${fp.split("/").at(-1)}:${lineno}`);
 function codeBlock(st, language, { loc, wrap }={}) {
  const code = document.createElement("code");
  // plaintext renders like a terminal print, otherwise render with syntax highlighting
@@ -547,11 +560,7 @@ function codeBlock(st, language, { loc, wrap }={}) {
  code.className = "hljs";
  const ret = document.createElement("pre");
  if (wrap) ret.className = "wrap";
-  if (loc != null) {
-    const link = ret.appendChild(document.createElement("a"));
-    link.href = "vscode://file/"+loc.join(":");
-    link.textContent = `${loc[0].split("/").at(-1)}:${loc[1]}`+"\n\n";
-  }
+  if (loc != null) ret.appendChild(pathLink(loc[0], loc[1]).style("margin-bottom", "4px").node());
  ret.appendChild(code);
  return ret;
 }
@@ -751,6 +760,15 @@ async function main() {
  // ** right sidebar code blocks
  const codeElement = codeBlock(ret[currentRewrite].uop, "python", { wrap:false });
  metadata.replaceChildren(toggleLabel, codeBlock(step.code_line, "python", { loc:step.loc, wrap:true }), codeElement);
+  if (step.trace) {
+    const trace = d3.create("pre").append("code").classed("hljs", true);
+    for (let i=step.trace.length-1; i>=0; i--) {
+      const [fp, lineno, fn, code] = step.trace[i];
+      trace.append("div").style("margin-bottom", "2px").style("display","flex").text(fn+" ").append(() => pathLink(fp, lineno).node());
+      trace.append("div").html(hljs.highlight(code, { language: "python" }).value).style("margin-bottom", "1ex");
+    }
+    metadata.insertBefore(trace.node().parentNode, codeElement);
+  }
  // ** rewrite steps
  if (step.match_count >= 1) {
    const rewriteList = metadata.appendChild(document.createElement("div"));
@@ -7,7 +7,8 @@ from http.server import BaseHTTPRequestHandler
 from urllib.parse import parse_qs, urlparse
 from typing import Any, TypedDict, TypeVar, Generator, Callable
 from tinygrad.helpers import colored, getenv, tqdm, unwrap, word_wrap, TRACEMETA, ProfileEvent, ProfileRangeEvent, TracingKey, ProfilePointEvent, temp
-from tinygrad.uop.ops import TrackedGraphRewrite, RewriteTrace, UOp, Ops, printable, GroupOp, srender, sint, sym_infer, range_str, pyrender
+from tinygrad.helpers import printable
+from tinygrad.uop.ops import TrackedGraphRewrite, RewriteTrace, UOp, Ops, GroupOp, srender, sint, sym_infer, range_str, pyrender
 from tinygrad.uop.ops import print_uops, range_start, multirange_str
 from tinygrad.device import ProfileDeviceEvent, ProfileGraphEvent, ProfileGraphEntry, Device
 from tinygrad.renderer import ProgramSpec
@@ -30,7 +31,7 @@ ref_map:dict[Any, int] = {}
 def get_rewrites(t:RewriteTrace) -> list[dict]:
  ret = []
  for i,(k,v) in enumerate(zip(t.keys, t.rewrites)):
-    steps = [{"name":s.name, "loc":s.loc, "match_count":len(s.matches), "code_line":printable(s.loc),
+    steps = [{"name":s.name, "loc":s.loc, "match_count":len(s.matches), "code_line":printable(s.loc), "trace":k.tb if j == 0 else None,
              "query":f"/ctxs?ctx={i}&idx={j}", "depth":s.depth} for j,s in enumerate(v)]
    if isinstance(k.ret, ProgramSpec):
      steps.append({"name":"View UOp List", "query":f"/render?ctx={i}&fmt=uops", "depth":0})
@@ -79,7 +80,7 @@ def uop_to_json(x:UOp) -> dict[int, dict]:
    try:
      if len(rngs:=u.ranges):
        label += f"\n({multirange_str(rngs, color=True)})"
-      if u.op not in {Ops.BUFFER, Ops.KERNEL, Ops.ASSIGN, Ops.COPY, Ops.SINK, *GroupOp.Buffer} and u._shape is not None:
+      if u._shape is not None:
        label += f"\n{shape_to_str(u.shape)}"
      if u.op in {Ops.INDEX, Ops.BUFFERIZE}:
        label += f"\n{u.render()}"
@@ -150,17 +151,22 @@ def timeline_layout(dev_events:list[tuple[int, int, float, DevEvent]], start_ts:
  for st,et,dur,e in dev_events:
    if isinstance(e, ProfilePointEvent) and e.name == "exec": exec_points[e.arg["name"]] = e
    if dur == 0: continue
-    name, info, key = e.name, None, None
+    name, fmt, key = e.name, [], None
    if (ref:=ref_map.get(name)) is not None:
      name = ctxs[ref]["name"]
      if isinstance(p:=trace.keys[ref].ret, ProgramSpec) and (ei:=exec_points.get(p.name)) is not None:
-        info = f"{sym_infer(p.estimates.ops, ei.arg['var_vals'])/(t:=dur*1e3):.2f} GFLOPS {sym_infer(p.estimates.mem, ei.arg['var_vals'])/t:4.1f}"+ \
-               f"|{sym_infer(p.estimates.lds,ei.arg['var_vals'])/t:.1f} GB/s\n{[str(m) for m in (ei.arg['metadata'] or ())]}"
+        flops = sym_infer(p.estimates.ops, var_vals:=ei.arg['var_vals'])/(t:=dur*1e-6)
+        membw, ldsbw = sym_infer(p.estimates.mem, var_vals)/t, sym_infer(p.estimates.lds, var_vals)/t
+        fmt = [f"{flops*1e-9:.0f} GFLOPS" if flops < 1e14 else f"{flops*1e-12:.0f} TFLOPS",
+              (f"{membw*1e-9:.0f} GB/s" if membw < 1e13 else f"{membw*1e-12:.0f} TB/s")+" mem",
+              (f"{ldsbw*1e-9:.0f} GB/s" if ldsbw < 1e15 else f"{ldsbw*1e-12:.0f} TB/s")+" lds"]
+        if (metadata_str:=",".join([str(m) for m in (ei.arg['metadata'] or ())])): fmt.append(metadata_str)
+        if isinstance(e, ProfileGraphEntry): fmt.append("(batched)")
        key = ei.key
    elif isinstance(e.name, TracingKey):
      name = e.name.display_name
      ref = next((v for k in e.name.keys if (v:=ref_map.get(k)) is not None), None)
-    events.append(struct.pack("<IIIIfI", enum_str(name, scache), option(ref), option(key), st-start_ts, dur, enum_str(info or "", scache)))
+    events.append(struct.pack("<IIIIfI", enum_str(name, scache), option(ref), option(key), st-start_ts, dur, enum_str("\n".join(fmt), scache)))
  return struct.pack("<BI", 0, len(events))+b"".join(events) if events else None

 def encode_mem_free(key:int, ts:int, execs:list[ProfilePointEvent], scache:dict) -> bytes:
@@ -212,12 +218,23 @@ def load_sqtt(profile:list[ProfileEvent]) -> None:
    if (r:=ref_map.get(name)): name = ctxs[r]["name"]
    steps.append({"name":name, "depth":0, "query":f"/render?ctx={len(ctxs)}&step={len(steps)}&fmt=counters",
                  "data":{"src":trace.keys[r].ret.src if r else name, "lang":"cpp"}})
+
+    # Idle:     The total time gap between the completion of previous instruction and the beginning of the current instruction.
+    #           The idle time can be caused by:
+    #             * Arbiter loss
+    #             * Source or destination register dependency
+    #             * Instruction cache miss
+    # Stall:    The total number of cycles the hardware pipe couldn't issue an instruction.
+    # Duration: Total latency in cycles, defined as "Stall time + Issue time" for gfx9 or "Stall time + Execute time" for gfx10+.
    for w in waves:
-      rows = [(e.inst, e.time, e.time-(w.insts[i-1].time if i else 0), e.dur, e.stall, str(e.typ).split("_")[-1]) for i,e in enumerate(w.insts)]
-      summary = [{"label":"Total Cycles", "value":w.insts[-1].time-w.insts[0].time if w.insts else 0}, {"label":"CU", "value":w.cu},
+      rows, prev_instr = [], w.begin_time
+      for i,e in enumerate(w.insts):
+        rows.append((e.inst, e.time, max(0, e.time-prev_instr), e.dur, e.stall, str(e.typ).split("_")[-1]))
+        prev_instr = max(prev_instr, e.time + e.dur)
+      summary = [{"label":"Total Cycles", "value":w.end_time-w.begin_time}, {"label":"CU", "value":w.cu},
                 {"label":"SIMD", "value":w.simd}]
      steps.append({"name":f"Wave {w.wave_id}", "depth":1, "query":f"/render?ctx={len(ctxs)}&step={len(steps)}&fmt=counters",
-                    "data":{"rows":rows, "cols":["Instruction", "Clk", "Wait", "Duration", "Stall", "Type"], "summary":summary}})
+                    "data":{"rows":rows, "cols":["Instruction", "Clk", "Idle", "Duration", "Stall", "Type"], "summary":summary}})
  ctxs.append({"name":"Counters", "steps":steps})

 def get_profile(profile:list[ProfileEvent]) -> bytes|None:
Author	SHA1	Message	Date
George HotzandGitHub	86d1e42ed8	Merge branch 'master' into algebraic_upat	2025-11-13 09:21:05 -08:00
George HotzandGitHub	faf68c03a8	more mi350x matmul work (#13138 ) * more mi350x matmul work * broken compute	2025-11-13 09:09:28 -08:00
George HotzandGitHub	5c13504bc1	Merge branch 'master' into algebraic_upat	2025-11-13 09:07:54 -08:00
Ayman Jabr GitHub George Hotz	256f81bb02	Fix tracemeta 0 (#13049 ) * chore: tclesius branch resolved * fix: indentation --------- Co-authored-by: George Hotz <[email protected]>	2025-11-13 09:07:11 -08:00
alpharushandGitHub	7e0aaadecd	feat: add repro command to summary (#10930 )	2025-11-13 08:52:27 -08:00
nimlgenandGitHub	6be86dde17	nv: add timeout when repsonding to rpc (#13260 )	2025-11-14 00:42:21 +08:00
nimlgenandGitHub	f9b7586e08	roc: fix blob gc (#13256 )	2025-11-13 23:38:35 +08:00
George HotzandGitHub	263b724143	one cache and bump it (#13258 )	2025-11-13 07:33:31 -08:00
George HotzandGitHub	5efa727b83	move _pool to MovementMixins (#13257 )	2025-11-13 07:28:52 -08:00
geohot	bcdfc109b5	hotfix: disable flaky test	2025-11-13 06:19:28 -08:00
qazalandGitHub	006dea4c3e	roc: only save instruction execs (#13254 )	2025-11-13 21:28:40 +08:00
nimlgenandGitHub	f9586b38ba	system: pci mask and val (#13251 )	2025-11-13 20:44:58 +08:00
George HotzandGitHub	7316da3253	new readme (#13250 ) * new readme * update	2025-11-13 00:48:28 -08:00
geohot	17aa3379e9	hotfix: improve self_tokenize	2025-11-13 00:18:57 -08:00
chenyuandGitHub	4e5a9132e7	JIT_BATCH_SIZE=0 in compile3 (#13245 ) fixed some enqueue time	2025-11-12 23:12:45 -05:00
wozeparrotandGitHub	759557f633	feat: move tk tests to testextra (#13242 )	2025-11-12 17:06:53 -08:00
chenyuandGitHub	3f939f3d3c	update pm_simplify_valid (#13241 ) * update pm_simplify_valid fixed openpilot conv regression * IMAGE training is broken	2025-11-12 19:40:02 -05:00
chenyuandGitHub	f9851a852f	minor update to uop_given_valid [pr] (#13243 ) split from #13241	2025-11-12 19:03:18 -05:00
qazalandGitHub	fe2876a6d8	hotfix: second GB/s in viz (#13240 )	2025-11-13 07:14:27 +08:00
George HotzandGitHub	a23dea202b	actually make AMD_LLVM not default (#13238 )	2025-11-12 15:07:23 -08:00
George HotzandGitHub	ab9fa964d8	DISABLE_COMPILER_CACHE -> CCACHE (#13234 ) * DISABLE_COMPILER_CACHE -> CCACHE * Fix cachekey assignment in Compiler constructor	2025-11-12 15:07:09 -08:00
qazalandGitHub	be2e24cb25	roc: requires sudo to install (#13237 )	2025-11-12 16:59:22 -05:00
geohot	8f1f195b6d	hotfix: no hexdump for usbgpu patch.py	2025-11-12 12:05:37 -08:00
nimlgenandGitHub	9a53fcbde4	amd: sqtt on rdna3.5 (#13233 )	2025-11-13 03:30:42 +08:00
George HotzandGitHub	13f10a31dc	AMD_LLVM default off (#13232 )	2025-11-12 11:06:33 -08:00
qazalandGitHub	8b26cf2b3d	sqtt: update rcp timing test (#13231 ) * sqtt: assert correct output in timing test * found why	2025-11-13 02:01:54 +08:00
Jan AkhremchikandGitHub	bc8e537423	Add NONZERO op to onnx backend (#13211 )	2025-11-12 08:55:51 -08:00
nimlgenandGitHub	af17e07251	viz: sqtt touchups (#13228 ) * viz: sqtt touchups * revert * matches	2025-11-12 22:40:37 +08:00
qazalandGitHub	7a6853fa40	viz: show python callstack in the first graph (#13218 )	2025-11-12 20:52:28 +08:00
nimlgenandGitHub	82eb63d3ad	qcom: auto switch idle timer when profiling (#13230 ) * qcom: auto switch idle timer when profiling * fi	2025-11-12 20:31:24 +08:00
nimlgenandGitHub	fcd8d0751a	test_timing for hip (#13229 )	2025-11-12 20:28:58 +08:00
qazalandGitHub	74b9d33acb	viz: direct link to program source (#13227 )	2025-11-12 16:27:13 +08:00
wozeparrotandGitHub	371c1f2355	tk: move tiles to class (#13224 )	2025-11-11 21:53:46 -08:00
sirhcmandGitHub	41a098a82d	In-tree autogen: libc.py (#13217 ) * checkout changes from autogen branch * parents * pylint happy * move sys to system in helpers.py * typo * typo	2025-11-11 19:13:48 -08:00
wozeparrotandGitHub	222bb12ddf	tk softmax (#13205 )	2025-11-11 15:13:16 -08:00
wozeparrotandGitHub	787f0070ed	feat: don't use output reg as local reduce reg (#13203 )	2025-11-11 14:35:16 -08:00
chenyuandGitHub	ece1415def	clean up image_dot and image_conv2d (#13222 ) * clean up image_dot and image_conv2d * those are fine * interesting	2025-11-11 15:53:03 -05:00
nimlgenandGitHub	2f0ea29b34	qcom: 48bit timestamps (#13214 ) * qcom: 48bit timestamps * f * lol * fix	2025-11-12 04:14:33 +08:00
qazalandGitHub	bc55bc4849	cleanup test_viz profiler tests (#13221 )	2025-11-12 03:46:48 +08:00
chenyuandGitHub	23b90945c3	add a benchmark for openpilot vision with DEBUG=2 (#13219 ) see per kernel speed, also disable the jobs for 0.9.9	2025-11-11 14:41:52 -05:00
George HotzandGitHub	c2075f3613	gc disable during big rewrites (#13215 ) * gc disable during big rewrites * cleaner with helper	2025-11-11 10:30:47 -08:00
Roelof van DijkandGitHub	e59313da08	migrate pytest and ruff (#13216 )	2025-11-11 13:27:51 -05:00
Gaétan LepageandGitHub	6fd7ce3832	migrate to pyproject.toml (#13189 ) * migrate to pyproject.toml * move mypy config to pyproject.toml	2025-11-11 09:09:27 -08:00
qazalandGitHub	8002921a04	viz: improve the program run tooltip (#13212 ) * add tflops to tooltip format * show if the run was batched	2025-11-12 00:56:03 +08:00
qazalandGitHub	f91e366a17	viz: display the graph layout recursion error (#13194 ) * viz: display the graph layout recursion error * share styles * +min-width * same thing * inline the append	2025-11-11 15:25:12 +08:00
wozeparrotandGitHub	73497af4c0	clean: use np for allclose (#13204 )	2025-11-10 23:02:43 -08:00
George HotzandGitHub	a6360fd94d	store can have shape (#13202 ) * store can have shape * _shape	2025-11-10 22:16:47 -08:00
b1tg GitHub chenyu	f3692b7406	clean up hip renderer (#13063 ) * clean up hip renderer * ocml --------- Co-authored-by: chenyu <[email protected]>	2025-11-11 00:44:24 -05:00
chenyuandGitHub	22b8579234	one last regressed dm kernel (#13201 )	2025-11-10 23:30:52 -05:00
chenyuandGitHub	58b7e4fab3	GROUPTOP heuristic on more axes (#13206 ) fixed dm speed	2025-11-10 23:30:37 -05:00
geohot	6538935441	write tests for algebraic UPat	2025-10-05 08:21:38 +08:00