Bump CUDA Compiler JLL, test different compilers.

.buildkite/pipeline.yml

@@ -1,5 +1,7 @@
 steps:
+  # first, test supported Julia versions
   - group: ":julia: Julia"
+    key: "julia"
     steps:
       - label: "Julia {{matrix.julia}}"
         plugins:
@@ -10,7 +12,7 @@ steps:
         agents:
           queue: "cuda"
           gpu: "a100"
-        timeout_in_minutes: 15
+        timeout_in_minutes: 25
         matrix:
           setup:
             julia:
@@ -24,6 +26,37 @@ steps:
             build.message !~ /\[skip tests\]/ &&
             build.message !~ /\[skip julia\]/
 
+  # then, test supported CUDA compilers
+  - group: "CUDA"
+    key: "cuda"
+    depends_on: "julia"
+    steps:
+      - label: "CUDA {{matrix.cuda}}"
+        plugins:
+          - JuliaCI/julia#v1:
+              version: "1.12"
+          - JuliaCI/julia-test#v1:
+              coverage: false
+        agents:
+          queue: "cuda"
+          gpu: "a100"
+        if: |
+          build.message =~ /\[only tests\]/ ||
+          build.message =~ /\[only cuda\]/ ||
+          build.message !~ /\[only/ && !build.pull_request.draft &&
+            build.message !~ /\[skip tests\]/ &&
+            build.message !~ /\[skip cuda\]/
+        timeout_in_minutes: 25
+        matrix:
+          setup:
+            cuda:
+              - "13.4"
+              - "13.3"
+              - "13.2"
+              - "13.1"
+        commands: |
+          echo -e "[CUDA_Compiler_jll]\nversion = \"{{matrix.cuda}}\"" >LocalPreferences.toml
+
   - wait: ~
     continue_on_failure: true
 

LocalPreferences.toml

@@ -22,3 +22,11 @@
 
 # maximum disk cache size in bytes. defaults to 1 GiB.
 #cache_size_bytes = 1073741824
+
+
+[CUDA_Compiler_jll]
+# which CUDA compiler to use. it should not be necessary to set this preference.
+#version = "13.4"
+
+# whether to use a local CUDA compiler.
+#local = "false"

Project.toml

@@ -25,6 +25,9 @@ Random = "9a3f8284-a2c9-5f02-9a11-845980a1fd5c"
 SHA = "ea8e919c-243c-51af-8825-aaa63cd721ce"
 Scratch = "6c6a2e73-6563-6170-7368-637461726353"
 
+[sources]
+CUDACore = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "CUDACore"}
+
 [weakdeps]
 DLFP8Types = "f4c16678-4a16-415b-82ef-ed337c5d6c7c"
 Microfloats = "31c70f10-a750-4521-b13c-797315ae2933"
@@ -37,7 +40,7 @@ MicrofloatsExt = "Microfloats"
 Adapt = "4.4"
 BFloat16s = "0.6"
 CUDACore = "6.1"
-CUDA_Compiler_jll = "0.4"
+CUDA_Compiler_jll = "0.5"
 CUDA_Tile_jll = "13.3"
 CompilerCaching = "0.3.1, 0.4"
 DLFP8Types = "0.1.0"

docs/Project.toml

@@ -1,11 +1,30 @@
 [deps]
 CUDA = "052768ef-5323-5732-b1bb-66c8b64840ba"
+CUDACore = "bd0ed864-bdfe-4181-a5ed-ce625a5fdea2"
+CUDATools = "9ec180c6-1c07-47c7-9e6e-ebefa4d1f6d0"
+CUPTI = "9e67e8f6-ba02-4b6c-a7db-3b11ae1e7ab7"
 DLFP8Types = "f4c16678-4a16-415b-82ef-ed337c5d6c7c"
 Documenter = "e30172f5-a6a5-5a46-863b-614d45cd2de4"
 Microfloats = "31c70f10-a750-4521-b13c-797315ae2933"
+NVML = "611af6d1-644e-4c5d-bd58-854d7d1254b9"
+cuBLAS = "182d3088-87b7-4494-8cad-fc6afaa545bc"
+cuFFT = "533571aa-0936-420e-b4be-9c66f5f626ca"
+cuRAND = "20fd9a0b-12d5-4c2f-a8af-7c34e9e60431"
+cuSOLVER = "887afef0-6a32-4de5-add4-7827692ba8fc"
+cuSPARSE = "b26da814-b3bc-49ef-b0ee-c816305aa060"
 cuTile = "0dea8319-8c4a-4662-a73d-20234d115b9a"
 
 [sources]
+CUDA = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main"}
+CUDACore = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "CUDACore"}
+CUDATools = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "CUDATools"}
+CUPTI = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/cupti"}
+NVML = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/nvml"}
+cuBLAS = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/cublas"}
+cuFFT = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/cufft"}
+cuRAND = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/curand"}
+cuSOLVER = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/cusolver"}
+cuSPARSE = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/cusparse"}
 cuTile = {path = ".."}
 
 [compat]

examples/Project.toml

@@ -1,11 +1,30 @@
 [deps]
 CUDA = "052768ef-5323-5732-b1bb-66c8b64840ba"
+CUDACore = "bd0ed864-bdfe-4181-a5ed-ce625a5fdea2"
+CUDATools = "9ec180c6-1c07-47c7-9e6e-ebefa4d1f6d0"
+CUPTI = "9e67e8f6-ba02-4b6c-a7db-3b11ae1e7ab7"
 FFTW = "7a1cc6ca-52ef-59f5-83cd-3a7055c09341"
+NVML = "611af6d1-644e-4c5d-bd58-854d7d1254b9"
 NVTX = "5da4648a-3479-48b8-97b9-01cb529c0a1f"
 Random = "9a3f8284-a2c9-5f02-9a11-845980a1fd5c"
+cuBLAS = "182d3088-87b7-4494-8cad-fc6afaa545bc"
+cuFFT = "533571aa-0936-420e-b4be-9c66f5f626ca"
+cuRAND = "20fd9a0b-12d5-4c2f-a8af-7c34e9e60431"
+cuSOLVER = "887afef0-6a32-4de5-add4-7827692ba8fc"
+cuSPARSE = "b26da814-b3bc-49ef-b0ee-c816305aa060"
 cuTile = "0dea8319-8c4a-4662-a73d-20234d115b9a"
 
 [sources]
+CUDA = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main"}
+CUDACore = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "CUDACore"}
+CUDATools = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "CUDATools"}
+CUPTI = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/cupti"}
+NVML = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/nvml"}
+cuBLAS = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/cublas"}
+cuFFT = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/cufft"}
+cuRAND = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/curand"}
+cuSOLVER = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/cusolver"}
+cuSPARSE = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/cusparse"}
 cuTile = {path = ".."}
 
 [compat]

src/launch.jl

@@ -164,13 +164,36 @@ end
 
 public TileCompilerTimeoutError
 
+"""
+    tileiras_available() -> Bool
+
+Whether a `tileiras` binary is available, either through the `tileiras`
+preference or through `CUDA_Compiler_jll`. The JLL lacks `tileiras` when its
+selected CUDA version predates the compiler (CUDA 13.2) or when no artifact is
+available for the platform.
+"""
+tileiras_available() =
+    tileiras_override !== nothing ||
+    (CUDA_Compiler_jll.is_available() && isdefined(CUDA_Compiler_jll, :tileiras))
+
+function check_tileiras_available()
+    tileiras_available() && return
+    error("The selected CUDA_Compiler_jll (CUDA $(CUDA_Compiler_jll.cuda_version)) " *
+          "does not provide `tileiras`; select CUDA 13.2 or newer, or set the " *
+          "`tileiras` preference to a local binary.")
+end
+
 """
     tileiras_path() -> String
 
 Path to the `tileiras` binary. Honors the `tileiras` preference when set,
 otherwise falls back to `CUDA_Compiler_jll.tileiras_path`.
 """
-tileiras_path() = something(tileiras_override, CUDA_Compiler_jll.tileiras_path)
+function tileiras_path()
+    tileiras_override === nothing || return tileiras_override
+    check_tileiras_available()
+    return CUDA_Compiler_jll.tileiras_path
+end
 
 """
     tileiras_root() -> String
@@ -190,6 +213,7 @@ Construct a Cmd to invoke `tileiras` with `args`, with `CUDA_ROOT` set
 to [`tileiras_root`](@ref).
 """
 function tileiras_cmd(args...)
+    check_tileiras_available()
     cmd = tileiras_override === nothing ?
         `$(CUDA_Compiler_jll.tileiras()) $args` :
         Cmd([tileiras_override, args...])
@@ -229,6 +253,15 @@ function discover_tileir_disassembler()
         return TileIRDisassembler(`$disasm`, "--print-debug-info",
                                   parse_tileiras_version(log))
     end
+    # Prefer the disassembler shipped next to `tileiras` (CUDA 13.4+), which is
+    # version-matched to the compiler and thus decodes any bytecode we emit.
+    if CUDA_Compiler_jll.is_available() && isdefined(CUDA_Compiler_jll, :tileirdisasm)
+        proc, log = run_and_collect(`$(CUDA_Compiler_jll.tileirdisasm()) --version`)
+        success(proc) || error("tileirdisasm --version failed with exit code " *
+                               "$(proc.exitcode):\n$log")
+        return TileIRDisassembler(CUDA_Compiler_jll.tileirdisasm(), "--print-debug-info",
+                                  parse_tileiras_version(log))
+    end
     CUDA_Tile_jll.is_available() || error("CUDA_Tile_jll is not available")
     translate = `$(CUDA_Tile_jll.cuda_tile_translate()) --cudatilebc-to-mlir`
     return TileIRDisassembler(translate, "--mlir-print-debuginfo",
@@ -391,9 +424,7 @@ const tileir_toolchain_cache = LazyInitialized{TileIRToolchain}()
 function discover_tileir_toolchain()
     bytecode_version_override === nothing ||
         validate_bytecode_version(bytecode_version_override)
-    if tileiras_override === nothing && !CUDA_Compiler_jll.is_available()
-        error("CUDA_Compiler_jll is not available and no `tileiras` preference is set")
-    end
+    check_tileiras_available()
 
     proc, identity = run_and_collect(tileiras_cmd("--version"))
     success(proc) || error("tileiras --version failed with exit code " *

src/precompile.jl

@@ -40,6 +40,9 @@ import REPL
     # `tileiras` → CUBIN without needing a CUDA context. The runtime path
     # tacks `link` on top to load the CUBIN onto the GPU.
     function precompile_kernel(@nospecialize(f), @nospecialize(tt))
+        # Without `tileiras` (e.g. a pre-13.2 CUDA_Compiler_jll) kernels cannot
+        # compile; still allow the package itself to precompile.
+        tileiras_available() || return
         argtypes, const_argtypes = unwrap_argtypes(f, tt)
         bv = bytecode_version()
         for sm_arch in [v"8.0", v"8.6", v"8.7", v"8.9",

test/Project.toml

@@ -1,5 +1,8 @@
 [deps]
 CUDA = "052768ef-5323-5732-b1bb-66c8b64840ba"
+CUDACore = "bd0ed864-bdfe-4181-a5ed-ce625a5fdea2"
+CUDATools = "9ec180c6-1c07-47c7-9e6e-ebefa4d1f6d0"
+CUPTI = "9e67e8f6-ba02-4b6c-a7db-3b11ae1e7ab7"
 DLFP8Types = "f4c16678-4a16-415b-82ef-ed337c5d6c7c"
 FFTW = "7a1cc6ca-52ef-59f5-83cd-3a7055c09341"
 FileCheck = "4e644321-382b-4b05-b0b6-5d23c3d944fb"
@@ -15,8 +18,25 @@ Printf = "de0858da-6303-5e67-8744-51eddeeeb8d7"
 Random = "9a3f8284-a2c9-5f02-9a11-845980a1fd5c"
 SHA = "ea8e919c-243c-51af-8825-aaa63cd721ce"
 Test = "8dfed614-e22c-5e08-85e1-65c5234f0b40"
+cuBLAS = "182d3088-87b7-4494-8cad-fc6afaa545bc"
+cuFFT = "533571aa-0936-420e-b4be-9c66f5f626ca"
+cuRAND = "20fd9a0b-12d5-4c2f-a8af-7c34e9e60431"
+cuSOLVER = "887afef0-6a32-4de5-add4-7827692ba8fc"
+cuSPARSE = "b26da814-b3bc-49ef-b0ee-c816305aa060"
 cuTile = "0dea8319-8c4a-4662-a73d-20234d115b9a"
 
+[sources]
+CUDA = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main"}
+CUDACore = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "CUDACore"}
+CUDATools = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "CUDATools"}
+CUPTI = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/cupti"}
+NVML = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/nvml"}
+cuBLAS = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/cublas"}
+cuFFT = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/cufft"}
+cuRAND = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/curand"}
+cuSOLVER = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/cusolver"}
+cuSPARSE = {url = "https://github.com/JuliaGPU/CUDA.jl", rev = "main", subdir = "lib/cusparse"}
+
 [compat]
 CUDA = "6.1"
 FFTW = "1.10.0"

test/codegen/boundscheck_views.jl

@@ -95,7 +95,7 @@ has_padding(f, argtypes; bytecode_version) =
         @test has_padding(f, two_dim; bytecode_version=v"13.3")
     end
 
-    if ct.bytecode_version() >= v"13.4"
+    if ct.bytecode_version() >= v"13.4" && ct.tileir_disassembler_version() >= v"13.4"
         # An explicit false is the only way to select the v13.4 `inbounds`
         # encoding. It also drops padding, which Tile IR rejects on an
         # unchecked view.