python3Packages.tinygrad: 0.11.0 -> 0.12.0 (#479427)

pkgs/development/python-modules/tinygrad/default.nix

@@ -17,35 +17,43 @@
   setuptools,
 
   # optional-dependencies
-  llvmlite,
-  triton,
+  # arm
   unicorn,
+  # triton
+  triton,
+  # testing_minimal
+  hypothesis,
+  numpy,
+  pytest-xdist,
+  torch,
+  z3-solver,
+  # testing_unit
+  ggml-python,
+  openai,
+  safetensors,
+  tabulate,
+  tqdm,
+  # testing
+  blobfile,
+  boto3,
+  bottle,
+  capstone,
+  librosa,
+  networkx,
+  nibabel,
+  onnx,
+  onnxruntime,
+  opencv4,
+  pandas,
+  pillow,
+  pycocotools,
+  sentencepiece,
+  tiktoken,
+  transformers,
 
   # tests
   pytestCheckHook,
   writableTmpDirAsHomeHook,
-  blobfile,
-  bottle,
-  capstone,
-  clang,
-  hexdump,
-  hypothesis,
-  jax,
-  librosa,
-  ml-dtypes,
-  networkx,
-  numpy,
-  onnx,
-  onnxruntime,
-  pillow,
-  pytest-xdist,
-  safetensors,
-  sentencepiece,
-  tiktoken,
-  torch,
-  tqdm,
-  transformers,
-  z3-solver,
 
   # passthru
   tinygrad,
@@ -56,118 +64,124 @@
 
 buildPythonPackage (finalAttrs: {
   pname = "tinygrad";
-  version = "0.11.0";
+  version = "0.12.0";
   pyproject = true;
 
   src = fetchFromGitHub {
     owner = "tinygrad";
     repo = "tinygrad";
     tag = "v${finalAttrs.version}";
-    hash = "sha256-VG2rhkiwPFN3JYSBbqrwCdqhdGE8GY6oEatMSCydhw8=";
+    hash = "sha256-Lied19C1sAbislr2WznnCZEmOn5PA0OzMg2KOdWOYkA=";
   };
 
-  patches = [
-    (replaceVars ./fix-dlopen-cuda.patch {
-      inherit (addDriverRunpath) driverLink;
-      libnvrtc =
-        if cudaSupport then
-          "${lib.getLib cudaPackages.cuda_nvrtc}/lib/libnvrtc.so"
-        else
-          "Please import nixpkgs with `config.cudaSupport = true`";
-    })
-  ];
+  patches =
+    let
+      libExtension = stdenv.hostPlatform.extensions.sharedLibrary;
+    in
+    [
+      (replaceVars ./patch-deps-paths.patch {
+        libllvm = "${lib.getLib llvmPackages.llvm}/lib/libLLVM${libExtension}";
+        libclang = "${lib.getLib llvmPackages.libclang}/lib/libclang${libExtension}";
 
-  postPatch =
-    # Patch `clang` directly in the source file
-    # Use the unwrapped variant to enable the "native" features currently unavailable in the sandbox
-    ''
-      substituteInPlace tinygrad/runtime/ops_cpu.py \
-        --replace-fail "getenv(\"CC\", 'clang')" "'${lib.getExe llvmPackages.clang-unwrapped}'"
-    ''
-    + ''
-      substituteInPlace tinygrad/runtime/support/llvm.py \
-        --replace-fail "ctypes.util.find_library('LLVM')" "'${lib.getLib llvmPackages.llvm}/lib/libLLVM.so'"
-    ''
-    + lib.optionalString stdenv.hostPlatform.isLinux ''
-      substituteInPlace tinygrad/runtime/autogen/libc.py \
-        --replace-fail "ctypes.util.find_library('c')" "'${stdenv.cc.libc}/lib/libc.so.6'"
+        # Use the unwrapped variant to enable the "native" features currently unavailable in the sandbox
+        clang = lib.getExe llvmPackages.clang-unwrapped;
+      })
+    ]
+    ++ lib.optionals cudaSupport [
+      (replaceVars ./patch-cuda-paths.patch {
+        inherit (addDriverRunpath) driverLink;
+        cuda_nvrtc = lib.getLib cudaPackages.cuda_nvrtc;
 
-      substituteInPlace tinygrad/runtime/autogen/opencl.py \
-        --replace-fail "ctypes.util.find_library('OpenCL')" "'${ocl-icd}/lib/libOpenCL.so'"
-    ''
-    # test/test_tensor.py imports the PTX variable from the cuda_compiler.py file.
-    # This import leads to loading the libnvrtc.so library that is not substituted when cudaSupport = false.
-    # -> As a fix, we hardcode this variable to False
-    + lib.optionalString (!cudaSupport) ''
-      substituteInPlace test/test_tensor.py \
-        --replace-fail "from tinygrad.runtime.support.compiler_cuda import PTX" "PTX = False"
-    ''
-    # `cuda_fp16.h` and co. are needed at runtime to compile kernels
-    + lib.optionalString cudaSupport ''
-      substituteInPlace tinygrad/runtime/support/compiler_cuda.py \
-        --replace-fail \
-        '"-I/usr/local/cuda/include", "-I/usr/include", "-I/opt/cuda/include"' \
-        '"-I${lib.getDev cudaPackages.cuda_cudart}/include/"'
-    ''
-    + lib.optionalString rocmSupport ''
-      substituteInPlace tinygrad/runtime/autogen/hip.py \
-        --replace-fail "/opt/rocm/" "${rocmPackages.clr}/"
+        # `cuda_fp16.h` and co. are needed at runtime to compile kernels
+        cuda_cudart = lib.getInclude cudaPackages.cuda_cudart;
+      })
+    ]
+    ++ lib.optionals rocmSupport [
+      (replaceVars ./patch-rocm-paths.patch {
+        comgr = lib.getLib rocmPackages.rocm-comgr;
+        clr = lib.getLib rocmPackages.clr;
+        rocm-runtime = lib.getLib rocmPackages.rocm-runtime;
+        rocm-llvm-objdump = lib.getExe' rocmPackages.llvm.llvm "llvm-objdump";
+        llvm-objdump = lib.getExe' llvmPackages.llvm "llvm-objdump";
+        hipcc = lib.getExe' rocmPackages.hipcc "hipcc";
+      })
+    ];
 
-      substituteInPlace tinygrad/runtime/support/compiler_hip.py \
-        --replace-fail "/opt/rocm/include" "${rocmPackages.clr}/include"
+  postPatch = lib.optionalString stdenv.hostPlatform.isLinux ''
+    substituteInPlace tinygrad/runtime/autogen/opencl.py \
+      --replace-fail \
+        "dll = DLL('opencl', 'OpenCL')" \
+        "dll = DLL('opencl', '${lib.getLib ocl-icd}/lib/libOpenCL.so')"
 
-      substituteInPlace tinygrad/runtime/support/compiler_hip.py \
-        --replace-fail "/opt/rocm/llvm" "${rocmPackages.llvm.llvm}"
-
-      substituteInPlace tinygrad/runtime/autogen/comgr.py \
-        --replace-fail "/opt/rocm/" "${rocmPackages.rocm-comgr}/"
-    '';
+    substituteInPlace tinygrad/runtime/autogen/libc.py \
+      --replace-fail \
+        "dll = DLL('libc', 'c', use_errno=True)" \
+        "dll = DLL('libc', '${lib.getLib stdenv.cc.libc}/lib/libc.so.6', use_errno=True)"
+  '';
 
   __propagatedImpureHostDeps = lib.optional stdenv.hostPlatform.isDarwin "/usr/lib/libc.dylib";
 
   build-system = [ setuptools ];
 
-  optional-dependencies = {
-    llvm = [ llvmlite ];
+  optional-dependencies = lib.fix (self: {
     arm = [ unicorn ];
     triton = [ triton ];
-  };
+    testing_minimal = [
+      hypothesis
+      numpy
+      pytest-xdist
+      torch
+      z3-solver
+    ];
+    testing_unit = self.testing_minimal ++ [
+      ggml-python
+      openai
+      safetensors
+      tabulate
+      tqdm
+    ];
+    testing = self.testing_unit ++ [
+      blobfile
+      boto3
+      bottle
+      capstone
+      librosa
+      networkx
+      nibabel
+      onnx
+      onnxruntime
+      opencv4
+      pandas
+      pillow
+      pycocotools
+      sentencepiece
+      tiktoken
+      transformers
+    ];
+  });
 
   pythonImportsCheck = [
     "tinygrad"
+    "tinygrad.runtime.autogen.libclang"
   ]
   ++ lib.optionals cudaSupport [
+    "tinygrad.runtime.ops_cuda"
     "tinygrad.runtime.ops_nv"
+  ]
+  ++ lib.optionals rocmSupport [
+    "tinygrad.runtime.ops_amd"
+    "tinygrad.runtime.ops_hip"
+  ]
+  ++ lib.optionals stdenv.hostPlatform.isDarwin [
+    "tinygrad.runtime.ops_metal"
   ];
 
   nativeCheckInputs = [
+    llvmPackages.clang
     pytestCheckHook
     writableTmpDirAsHomeHook
-
-    blobfile
-    bottle
-    capstone
-    clang
-    hexdump
-    hypothesis
-    jax
-    librosa
-    ml-dtypes
-    networkx
-    numpy
-    onnx
-    onnxruntime
-    pillow
-    pytest-xdist
-    safetensors
-    sentencepiece
-    tiktoken
-    torch
-    tqdm
-    transformers
-    z3-solver
   ]
-  ++ networkx.optional-dependencies.extra;
+  ++ finalAttrs.passthru.optional-dependencies.testing;
 
   disabledTests = [
     # RuntimeError: Attempting to relocate against an undefined symbol 'fmaxf'
@@ -192,6 +206,7 @@ buildPythonPackage (finalAttrs: {
     "test_dataset_is_realized"
     "test_e2e_big"
     "test_fetch_small"
+    "test_hevc_parser"
     "test_huggingface_enet_safetensors"
     "test_index_mnist"
     "test_linear_mnist"
@@ -206,8 +221,13 @@ buildPythonPackage (finalAttrs: {
     "test_load_convnext"
     "test_load_enet"
     "test_load_enet_alt"
+    "test_load_gpt2_q4_1"
     "test_load_llama2bfloat"
     "test_load_resnet"
+    "test_load_sample_mxfp4"
+    "test_load_sample_q6_k"
+    "test_load_tinyllama_q4_0"
+    "test_load_tinyllama_q8_0"
     "test_mnist_val"
     "test_openpilot_model"
     "test_resnet"
@@ -220,12 +240,19 @@ buildPythonPackage (finalAttrs: {
     "test_transcribe_long_no_batch"
     "test_vgg7"
   ]
-  ++ lib.optionals (stdenv.hostPlatform.system == "aarch64-linux") [
+  ++ lib.optionals (stdenv.hostPlatform.isLinux && stdenv.hostPlatform.isAarch64) [
     # Fail with AssertionError
     "test_casts_from"
     "test_casts_to"
+    "test_float_cast_to_unsigned_underflow"
     "test_int8"
     "test_int8_to_uint16_negative"
+  ]
+  ++ lib.optionals stdenv.hostPlatform.isDarwin [
+    # Flaky (pass when running a smaller set of tests: tests/unit/*, but not within the full test suite)
+    # AttributeError: module 'tinygrad.runtime.autogen.libclang' has no attribute 'clang_parseTranslationUnit'
+    "test_gen_from_header"
+    "test_struct_ordering"
   ];
 
   disabledTestPaths = [
@@ -236,10 +263,30 @@ buildPythonPackage (finalAttrs: {
 
     # Files under this directory are not considered as tests by upstream and should be skipped
     "extra/"
+  ]
+  ++ lib.optionals stdenv.hostPlatform.isDarwin [
+    # urllib.error.URLError: <urlopen error [SSL: CERTIFICATE_VERIFY_FAILED]
+    # certificate verify failed: self-signed certificate in certificate chain (_ssl.c:1032)>
+    "test/models/test_whisper.py"
   ];
 
-  passthru.tests = {
-    withCuda = tinygrad.override { cudaSupport = true; };
+  __darwinAllowLocalNetworking = true;
+
+  passthru = {
+    tests = {
+      withCuda = tinygrad.override { cudaSupport = true; };
+      withRocm = tinygrad.override { rocmSupport = true; };
+    };
+
+    gpuCheck = tinygrad.overridePythonAttrs (old: {
+      requiredSystemFeatures = [ "cuda" ];
+
+      pytestFlags = (old.pytestFlags or [ ]) ++ [
+        # When running in parallel, with GPU support, some tests become flaky:
+        # RuntimeError: Wait timeout: 30000 ms! (the signal is not set to 153, but 151)
+        "--maxprocesses=1"
+      ];
+    });
   };
 
   meta = {
@@ -248,10 +295,5 @@ buildPythonPackage (finalAttrs: {
     changelog = "https://github.com/tinygrad/tinygrad/releases/tag/${finalAttrs.src.tag}";
     license = lib.licenses.mit;
     maintainers = with lib.maintainers; [ GaetanLepage ];
-    badPlatforms = [
-      # Fatal Python error: Aborted
-      # onnxruntime/capi/_pybind_state.py", line 32 in <module>
-      "aarch64-linux"
-    ];
   };
 })

pkgs/development/python-modules/tinygrad/fix-dlopen-cuda.patch

@@ -1,49 +0,0 @@
-diff --git a/tinygrad/runtime/autogen/cuda.py b/tinygrad/runtime/autogen/cuda.py
-index a30c8f53..e2078ff6 100644
---- a/tinygrad/runtime/autogen/cuda.py
-+++ b/tinygrad/runtime/autogen/cuda.py
-@@ -145,7 +145,19 @@ def char_pointer_cast(string, encoding='utf-8'):
- 
- 
- _libraries = {}
--_libraries['libcuda.so'] = ctypes.CDLL(ctypes.util.find_library('cuda'))
-+libcuda = None
-+try:
-+    libcuda = ctypes.CDLL('libcuda.so')
-+except OSError:
-+    pass
-+try:
-+    libcuda = ctypes.CDLL('@driverLink@/lib/libcuda.so')
-+except OSError:
-+    pass
-+if libcuda is None:
-+    raise RuntimeError(f"`libcuda.so` not found")
-+
-+_libraries['libcuda.so'] = libcuda
- 
- 
- cuuint32_t = ctypes.c_uint32
-diff --git a/tinygrad/runtime/autogen/nvrtc.py b/tinygrad/runtime/autogen/nvrtc.py
-index 6af74187..c5a6c6c4 100644
---- a/tinygrad/runtime/autogen/nvrtc.py
-+++ b/tinygrad/runtime/autogen/nvrtc.py
-@@ -10,7 +10,18 @@ import ctypes, ctypes.util
- 
- 
- _libraries = {}
--_libraries['libnvrtc.so'] = ctypes.CDLL(ctypes.util.find_library('nvrtc'))
-+libnvrtc = None
-+try:
-+    libnvrtc = ctypes.CDLL('libnvrtc.so')
-+except OSError:
-+    pass
-+try:
-+    libnvrtc = ctypes.CDLL('@libnvrtc@')
-+except OSError:
-+    pass
-+if libnvrtc is None:
-+    raise RuntimeError(f"`libnvrtc.so` not found")
-+_libraries['libnvrtc.so'] = libnvrtc
- def string_cast(char_pointer, encoding='utf-8', errors='strict'):
-     value = ctypes.cast(char_pointer, ctypes.c_char_p).value
-     if value is not None and encoding is not None:

pkgs/development/python-modules/tinygrad/patch-cuda-paths.patch

@@ -0,0 +1,39 @@
+diff --git a/tinygrad/runtime/autogen/cuda.py b/tinygrad/runtime/autogen/cuda.py
+index 061c8d73f..65ec3ffd7 100644
+--- a/tinygrad/runtime/autogen/cuda.py
++++ b/tinygrad/runtime/autogen/cuda.py
+@@ -1,7 +1,7 @@
+ # mypy: ignore-errors
+ import ctypes
+ from tinygrad.runtime.support.c import DLL, Struct, CEnum, _IO, _IOW, _IOR, _IOWR
+-dll = DLL('cuda', 'cuda')
++dll = DLL('cuda', '@driverLink@/lib/libcuda.so')
+ cuuint32_t = ctypes.c_uint32
+ cuuint64_t = ctypes.c_uint64
+ CUdeviceptr_v2 = ctypes.c_uint64
+diff --git a/tinygrad/runtime/autogen/nvrtc.py b/tinygrad/runtime/autogen/nvrtc.py
+index 88085c45b..90518d403 100644
+--- a/tinygrad/runtime/autogen/nvrtc.py
++++ b/tinygrad/runtime/autogen/nvrtc.py
+@@ -2,7 +2,7 @@
+ import ctypes
+ from tinygrad.runtime.support.c import DLL, Struct, CEnum, _IO, _IOW, _IOR, _IOWR
+ import sysconfig
+-dll = DLL('nvrtc', 'nvrtc', f'/usr/local/cuda/targets/{sysconfig.get_config_vars().get("MULTIARCH", "").rsplit("-", 1)[0]}/lib')
++dll = DLL('nvrtc','@cuda_nvrtc@/lib/libnvrtc.so')
+ nvrtcResult = CEnum(ctypes.c_uint32)
+ NVRTC_SUCCESS = nvrtcResult.define('NVRTC_SUCCESS', 0)
+ NVRTC_ERROR_OUT_OF_MEMORY = nvrtcResult.define('NVRTC_ERROR_OUT_OF_MEMORY', 1)
+diff --git a/tinygrad/runtime/support/compiler_cuda.py b/tinygrad/runtime/support/compiler_cuda.py
+index 8f71a9255..fdbf01bad 100644
+--- a/tinygrad/runtime/support/compiler_cuda.py
++++ b/tinygrad/runtime/support/compiler_cuda.py
+@@ -43,7 +43,7 @@ def cuda_disassemble(lib:bytes, arch:str):
+ class CUDACompiler(Compiler):
+   def __init__(self, arch:str, cache_key:str="cuda"):
+     self.arch, self.compile_options = arch, [f'--gpu-architecture={arch}']
+-    self.compile_options += [f"-I{CUDA_PATH}/include"] if CUDA_PATH else ["-I/usr/local/cuda/include", "-I/usr/include", "-I/opt/cuda/include"]
++    self.compile_options += ["-I@cuda_cudart@/include"]
+     nvrtc_check(nvrtc.nvrtcVersion((nvrtcMajor := ctypes.c_int()), (nvrtcMinor := ctypes.c_int())))
+     if (nvrtcMajor.value, nvrtcMinor.value) >= (12, 4): self.compile_options.append("--minimal")
+     super().__init__(f"compile_{cache_key}_{self.arch}")

pkgs/development/python-modules/tinygrad/patch-deps-paths.patch

@@ -0,0 +1,39 @@
+diff --git a/tinygrad/runtime/autogen/libclang.py b/tinygrad/runtime/autogen/libclang.py
+index dbf22b6f5..4a5a097dc 100644
+--- a/tinygrad/runtime/autogen/libclang.py
++++ b/tinygrad/runtime/autogen/libclang.py
+@@ -1,7 +1,7 @@
+ # mypy: ignore-errors
+ import ctypes
+ from tinygrad.runtime.support.c import DLL, Struct, CEnum, _IO, _IOW, _IOR, _IOWR
+-dll = DLL('libclang', ['clang-20', 'clang'])
++dll = DLL('libclang', '@libclang@')
+ CXIndex = ctypes.c_void_p
+ class struct_CXTargetInfoImpl(Struct): pass
+ CXTargetInfo = ctypes.POINTER(struct_CXTargetInfoImpl)
+diff --git a/tinygrad/runtime/autogen/llvm.py b/tinygrad/runtime/autogen/llvm.py
+index 7d05224e7..54204f563 100644
+--- a/tinygrad/runtime/autogen/llvm.py
++++ b/tinygrad/runtime/autogen/llvm.py
+@@ -2,7 +2,7 @@
+ import ctypes
+ from tinygrad.runtime.support.c import DLL, Struct, CEnum, _IO, _IOW, _IOR, _IOWR
+ from tinygrad.helpers import WIN, OSX
+-dll = DLL('llvm', 'C:\\Program Files\\LLVM\\bin\\LLVM-C.dll' if WIN else '/opt/homebrew/opt/llvm@20/lib/libLLVM.dylib' if OSX else ['LLVM', 'LLVM-21', 'LLVM-20', 'LLVM-19', 'LLVM-18', 'LLVM-17', 'LLVM-16', 'LLVM-15', 'LLVM-14'])
++dll = DLL('llvm', '@libllvm@')
+ intmax_t = ctypes.c_int64
+ try: (imaxabs:=dll.imaxabs).restype, imaxabs.argtypes = intmax_t, [intmax_t]
+ except AttributeError: pass
+diff --git a/tinygrad/runtime/support/compiler_cpu.py b/tinygrad/runtime/support/compiler_cpu.py
+index 8b11f3af8..d8190fac3 100644
+--- a/tinygrad/runtime/support/compiler_cpu.py
++++ b/tinygrad/runtime/support/compiler_cpu.py
+@@ -15,7 +15,7 @@ class ClangJITCompiler(Compiler):
+     arch = {'x86_64': '-march=native', 'AMD64': '-march=native', 'riscv64': '-march=rv64g'}.get(platform.machine(), "-mcpu=native")
+     args = [arch, f'--target={target}-none-unknown-elf', '-O2', '-fPIC', '-ffreestanding', '-fno-math-errno', '-nostdlib', '-fno-ident']
+     arch_args = ['-ffixed-x18'] if target == 'arm64' else []
+-    obj = subprocess.check_output([getenv("CC", 'clang'), '-c', '-x', 'c', *args, *arch_args, '-', '-o', '-'], input=src.encode('utf-8'))
++    obj = subprocess.check_output(["@clang@", '-c', '-x', 'c', *args, *arch_args, '-', '-o', '-'], input=src.encode('utf-8'))
+     return jit_loader(obj)
+ 
+   def disassemble(self, lib:bytes): return capstone_flatdump(lib)

pkgs/development/python-modules/tinygrad/patch-rocm-paths.patch

@@ -0,0 +1,68 @@
+diff --git a/tinygrad/runtime/autogen/comgr.py b/tinygrad/runtime/autogen/comgr.py
+index 9c642cef2..b42c7c9c8 100644
+--- a/tinygrad/runtime/autogen/comgr.py
++++ b/tinygrad/runtime/autogen/comgr.py
+@@ -2,7 +2,7 @@
+ import ctypes
+ from tinygrad.runtime.support.c import DLL, Struct, CEnum, _IO, _IOW, _IOR, _IOWR
+ import os
+-dll = DLL('comgr', [os.getenv('ROCM_PATH', '/opt/rocm')+'/lib/libamd_comgr.so', 'amd_comgr'])
++dll = DLL('comgr', '@comgr@/lib/libamd_comgr.so')
+ amd_comgr_status_s = CEnum(ctypes.c_uint32)
+ AMD_COMGR_STATUS_SUCCESS = amd_comgr_status_s.define('AMD_COMGR_STATUS_SUCCESS', 0)
+ AMD_COMGR_STATUS_ERROR = amd_comgr_status_s.define('AMD_COMGR_STATUS_ERROR', 1)
+diff --git a/tinygrad/runtime/autogen/hip.py b/tinygrad/runtime/autogen/hip.py
+index 1cccb1e1f..d57a4d90a 100644
+--- a/tinygrad/runtime/autogen/hip.py
++++ b/tinygrad/runtime/autogen/hip.py
+@@ -2,7 +2,7 @@
+ import ctypes
+ from tinygrad.runtime.support.c import DLL, Struct, CEnum, _IO, _IOW, _IOR, _IOWR
+ import os
+-dll = DLL('hip', os.getenv('ROCM_PATH', '/opt/rocm')+'/lib/libamdhip64.so')
++dll = DLL('hip', '@clr@/lib/libamdhip64.so')
+ hipError_t = CEnum(ctypes.c_uint32)
+ hipSuccess = hipError_t.define('hipSuccess', 0)
+ hipErrorInvalidValue = hipError_t.define('hipErrorInvalidValue', 1)
+diff --git a/tinygrad/runtime/autogen/hsa.py b/tinygrad/runtime/autogen/hsa.py
+index 16e568e21..9788d3705 100644
+--- a/tinygrad/runtime/autogen/hsa.py
++++ b/tinygrad/runtime/autogen/hsa.py
+@@ -2,7 +2,7 @@
+ import ctypes
+ from tinygrad.runtime.support.c import DLL, Struct, CEnum, _IO, _IOW, _IOR, _IOWR
+ import os
+-dll = DLL('hsa', [os.getenv('ROCM_PATH', '/opt/rocm')+'/lib/libhsa-runtime64.so', 'hsa-runtime64'])
++dll = DLL('hsa', '@rocm-runtime@/lib/libhsa-runtime64.so')
+ enum_SQ_RSRC_BUF_TYPE = CEnum(ctypes.c_uint32)
+ SQ_RSRC_BUF = enum_SQ_RSRC_BUF_TYPE.define('SQ_RSRC_BUF', 0)
+ SQ_RSRC_BUF_RSVD_1 = enum_SQ_RSRC_BUF_TYPE.define('SQ_RSRC_BUF_RSVD_1', 1)
+diff --git a/tinygrad/runtime/support/compiler_amd.py b/tinygrad/runtime/support/compiler_amd.py
+index e97aaefcd..ef55838a3 100644
+--- a/tinygrad/runtime/support/compiler_amd.py
++++ b/tinygrad/runtime/support/compiler_amd.py
+@@ -13,9 +13,9 @@ from tinygrad.runtime.support.compiler_cpu import LLVMCompiler
+ from tinygrad.helpers import OSX, to_char_p_p
+ 
+ def _find_llvm_objdump():
+-  if OSX: return '/opt/homebrew/opt/llvm/bin/llvm-objdump'
++  if OSX: return '@llvm-objdump@'
+   # Try ROCm path first, then versioned, then unversioned
+-  for p in ['/opt/rocm/llvm/bin/llvm-objdump', 'llvm-objdump-21', 'llvm-objdump-20', 'llvm-objdump']:
++  for p in ['@rocm-llvm-objdump@', '@llvm-objdump@']:
+     if shutil.which(p): return p
+   raise FileNotFoundError("llvm-objdump not found")
+ 
+@@ -107,9 +107,9 @@ class HIPCCCompiler(Compiler):
+         srcf.write(src.encode())
+         srcf.flush()
+ 
+-        subprocess.run(["hipcc", "-c", "-emit-llvm", "--cuda-device-only", "-O3", "-mcumode",
+-                        f"--offload-arch={self.arch}", "-I/opt/rocm/include/hip", "-o", bcf.name, srcf.name] + self.extra_options, check=True)
+-        subprocess.run(["hipcc", "-target", "amdgcn-amd-amdhsa", f"-mcpu={self.arch}",
++        subprocess.run(["@hipcc@", "-c", "-emit-llvm", "--cuda-device-only", "-O3", "-mcumode",
++                        f"--offload-arch={self.arch}", "-I@clr@/include/hip", "-o", bcf.name, srcf.name] + self.extra_options, check=True)
++        subprocess.run(["@hipcc@", "-target", "amdgcn-amd-amdhsa", f"-mcpu={self.arch}",
+                         "-O3", "-mllvm", "-amdgpu-internalize-symbols", "-c", "-o", libf.name, bcf.name] + self.extra_options, check=True)
+ 
+         return pathlib.Path(libf.name).read_bytes()