State of Julia's GPU ecosystem in 2026

By: Guillaume Dalle

Re-posted from: https://juliagpu.org/post/2026-08-03-gpu_ecosystem/index.html

A summary of the various packages making up Julia's GPU abilities, and how they interact.

The text of this post was written by Claude Sonnet 4.6, then reviewed and edited by Guillaume Dalle and other contributors. The initial structure and list of packages had been manually curated beforehand.

Julia's GPU ecosystem has grown into a rich, layered stack that spans everything from vendor-specific low-level wrappers to hardware-agnostic high-level abstractions. This post gives an overview of the major packages, organized by where they sit in that stack. The distinction between hardware-specific and hardware-agnostic packages is the key design principle: vendor-specific backends provide raw access to each GPU platform, while a shared set of abstractions lets library authors and users write code that is portable across all of them.

Hardware-specific

CUDA ecosystem

The CUDA ecosystem is the most mature part of Julia's GPU stack, built around NVIDIA hardware.

CUDA.jl is the primary interface for programming NVIDIA GPUs in Julia. It bundles a user-friendly array abstraction (CuArray), a compiler for writing CUDA kernels directly in Julia, and can be supplemented with wrappers for a broad set of CUDA libraries including cuBLAS.jl, cuSPARSE.jl, cuFFT.jl, cuSOLVER.jl, and cuDNN.jl. Most Julia users who only target NVIDIA hardware start here and never need to go deeper.

cuTile.jl exposes NVIDIA's tile-based programming model, available on Ampere and newer GPUs, through a high-level Julia interface to the Tile IR architecture. It can fuse complex operations into single kernels while supporting specialized numeric types such as FP8 and mixed-precision formats that are central to modern machine learning workloads. Whereas CUDA.jl covers the breadth of CUDA, cuTile.jl is the tool of choice when squeezing maximum throughput out of NVIDIA's latest tensor cores.

CUDSS.jl is a Julia wrapper for NVIDIA's cuDSS library, which provides GPU-accelerated sparse linear solvers. It exposes three factorization methods (LDU, LDLᵀ, and LLᵀ) and fills a gap left by the main CUDA.jl bundle, since cuDSS remains in preview and is shipped separately.

cuNumeric.jl wraps NVIDIA's cuPyNumeric C++ API to bring distributed, multi-GPU array computing to Julia. It provides an NDArray abstraction that supports standard array operations (e.g., broadcasting, matmul) and automatically partitions work across multi-GPU systems without intervention from the user.

Other vendors

Beyond NVIDIA, Julia has backends for every major GPU platform.

AMDGPU.jl brings AMD GPU computing to Julia through ROCm integration. It mirrors the structure of CUDA.jl by providing an array type (ROCArray), a kernel compiler, and library wrappers for AMD's graphics and compute hardware.

oneAPI.jl targets Intel GPUs and accelerators through Intel's oneAPI unified programming toolkit. It provides low-level Level Zero API wrappers, a oneArray type that integrates with Julia's array ecosystem, and oneMKL bindings for optimized linear algebra and sparse matrix operations.

Metal.jl enables GPU programming on macOS using Apple's Metal framework, targeting Apple Silicon. The package offers three levels of abstraction: high-level array operations via MtlArray, custom kernel programming, and direct Metal API access through ObjectiveC bindings. While still under active development with some known limitations, it allows Mac users to run GPU-accelerated Julia code without any external hardware.

Hardware-agnostic

Data types

The hardware-agnostic layer starts with array types and the utilities to move data between them.

GPUArrays.jl is the foundational package that defines the shared interface all Julia GPU array types implement. Rather than serving end users directly, it establishes the AbstractGPUArray contract—analogous to Julia's AbstractArray—that backend developers implement when building types like CuArray, ROCArray, or MtlArray. The repository also ships two companion sub-packages: GPUArraysCore.jl, which provides the minimal type hierarchy for packages that only need to check whether an array is on a GPU, and JLArrays.jl, a CPU-backed reference implementation used for testing.

Adapt.jl provides a mechanism for converting wrapper types to GPU-compatible formats while preserving their structure. Unlike convert(), the adapt(T, x) function knows how to unwrap and re-wrap types like Adjoint or NamedTuple around GPU arrays rather than discarding them. GPU libraries including CUDA.jl use Adapt.jl's extension hooks (adapt_structure and adapt_storage) to make data movement to the device transparent, which is why user-defined structs containing arrays typically only need a single Adapt.@adapt_structure annotation to become GPU-compatible.

Low-level kernels

Two packages provide the primitives for writing custom GPU kernels in a portable way.

KernelAbstractions.jl is the central abstraction layer for writing GPU kernels that run across multiple hardware backends. It provides a unified, minimal @kernel macro that compiles to NVIDIA CUDA, AMD ROCm, Intel oneAPI, Apple Metal, OpenCL and the CPU without any backend-specific rewrites. Most hardware-agnostic libraries in Julia -— including AcceleratedKernels.jl or Lava.jl -— build on top of it, making it the glue that holds the portable GPU stack together.

KernelIntrinsics.jl provides low-level memory access primitives and warp-level operations for GPU kernel authors who need fine-grained control beyond what KernelAbstractions.jl exposes. It covers memory fencing, warp shuffle and reduction operations, and vectorized memory access (see the package documentation for details on what these are), and does so across CUDA, ROCm, and Metal backends. The package is aimed at library developers rather than end users: it fills the gap between high-level kernel abstractions and the raw hardware intrinsics that performance-critical GPU code sometimes requires.

OpenCL.jl provides a comprehensive Julia interface to the OpenCL standard, which targets GPUs, FPGAs, DSPs, and multicore CPUs from a single API. The package supports both traditional OpenCL C kernels and native Julia functions compiled to SPIR-V, making it the most broadly portable of the hardware-specific backends. It is a practical choice when targeting hardware not covered by the other backends, or when writing code that needs to run on a wide variety of devices. Through PoCL, it also provides a way of running GPU kernels on the CPU.

Vulkan.jl wraps the Vulkan graphics and compute API, generating bindings automatically from the official Vulkan specification with minimal overhead over the underlying C interface. Where OpenCL.jl offers portability at the cost of abstraction, Vulkan provides explicit, low-overhead control over GPU resources. The package hasn't reached 1.0 yet but is maintained and considered stable. It serves as a low-level foundation for higher-level graphics and compute work in Julia, and is rather meant for developers.

Lava.jl is a Julia GPU backend that compiles Julia code to SPIR-V for execution via Vulkan, functioning as a unified compute, graphics, and ray tracing platform. It serves as a drop-in replacement for other GPU backends through the KernelAbstractions.jl and GPUArrays.jl interface, while additionally enabling graphics shaders and hardware-accelerated ray tracing written entirely in Julia rather than GLSL. The package supports cross-platform execution on NVIDIA, AMD, Intel, Apple, and software renderers.

High-level programming

Several packages build on lower-level primitives to provide ready-made parallel algorithms.

AcceleratedKernels.jl provides cross-architecture parallel algorithms—sorting, reduction, accumulation, and more—that compile from a single codebase to multithreaded CPUs, CUDA, ROCm, oneAPI, and Metal. It has utilities for setting the number of threads, the block size, or pre-allocate scratchspaces.

GemmKernels.jl is a flexible framework for crafting optimized General Matrix Multiplication (GEMM) kernels on NVIDIA GPUs. It decomposes GEMM into modular, customizable components—parameters, layouts, transforms, operators, and epilogues—that users can mix and match through Julia's multiple dispatch system. The package can be useful when the standard BLAS interface is too inflexible for a particular memory layout or numeric type.

KernelForge.jl is a pure Julia library of high-performance, portable GPU primitives including map-reduce, prefix scans, matrix-vector products, and sorting. It targets both NVIDIA and AMD hardware and aims for performance comparable to optimized C++ libraries, without requiring any non-Julia dependencies.

JACC.jl provides a simple vendor-neutral API for CPU and GPU computing. It targets HPC users familiar with C++ frameworks like Kokkos, RAJA, SYCL or TBB. Its array construction (zeros/ones/fill), parallel_for and parallel_reduce primitives deploy to NVIDIA, AMD, Apple or Intel GPUs using JuliaGPU's vendor-specific backends. They can also leverage CPU threads using Polyester.jl. Backend selection is done outside code using Preferences.jl mechanisms (e.g., LocalPreferences.toml). The package is well-suited for HPC prototyping: developers can write and test kernels on a laptop CPU or GPU and then deploy them to multi-GPU supercomputer nodes without changing any application code. Users of the default APIs, do not need prior CPU/GPU programming knowledge to parallelize their codes, but JACC.jl provides low-level performance APIs (e.g., blocks, threads, async, shared memory, stream, multi-GPU, etc.) for hardware-specific optimizations.

Strided.jl provides a vendor-neutral API for writing map– or mapreduce– kernels over input arrays with varying strides. This allows for writing operations that fuse (strided) views and permutedims operations with the following kernel calls. StridedViews.jl represents lazy views with arbitrary strides over any subtype of DenseArray.

MatrixAlgebraKit.jl provides a high-level interface to linear algebra routines provided by the various GPU vendors. It features a unified way of accessing these kernels that exposes access to more in-place operations than LinearAlgebra.jl, as well as compatibility with the various automatic differentiation libraries.

Vendor detection and translation

As the number of backends grows, tooling for selecting and migrating between them becomes important.

GPUSelect.jl automates GPU backend selection for KernelAbstractions.jl by detecting available hardware at runtime through driver libraries. It provides applications with a one-liner interface to load the appropriate backend—whether CUDA, AMDGPU, Metal, oneAPI, or Vulkan—without manual configuration. The package is designed for end-user applications rather than libraries, handling both detection and, when needed, automatic installation of the relevant backend.

GPUEnv.jl simplifies multi-backend development by automatically detecting available GPU hardware and creating temporary overlay environments containing only the relevant backend packages. Rather than permanently including all GPU dependencies in a project, it conditionally activates only the packages that match the host machine's hardware using lightweight probe functions. This keeps parent environments lean and fast to resolve.

Juliana.jl is a translation tool that automatically converts Julia code written for CUDA.jl into portable multi-backend code compatible with KernelAbstractions.jl. This allows GPU programs originally written for NVIDIA hardware to run on Intel, AMD, and Apple GPUs without manual rewriting. It is most useful for porting existing CUDA.jl codebases toward hardware-agnostic designs without starting from scratch.

Linear algebra

NextLA.jl is a hardware-agnostic package containing implementations of BLAS/LAPACK routines for dense linear algebra. It supports multiple number types and leverages multi-threading as well as GPU acceleration.

Tensor operations

For operations on multi-dimensional arrays expressed through index notation, several packages provide GPU-aware implementations.

Tullio.jl provides a macro that translates index notation into optimized array operations, spanning multi-threading, SIMD vectorization, and GPU kernels through KernelAbstractions.jl. It handles complex patterns including convolutions, reductions, and scatter/gather, and supports automatic differentiation for machine learning workflows. Because the same @tullio expression dispatches to the appropriate backend based on the input array type, existing GPU arrays from CUDA.jl or AMDGPU.jl benefit automatically.

TensorCast.jl enables reshaping, permuting, slicing, and reducing multi-dimensional arrays using an intuitive index notation that compiles down to Julia's native broadcasting and array operations. When given GPU arrays from CUDA.jl or other backends, broadcasting operations execute directly on the device, so the package integrates naturally into GPU workflows without requiring any special GPU-specific code paths. It is particularly useful for expressing data layout transformations that would otherwise require verbose combinations of reshape, permutedims, and dropdims.

OMEinsum.jl implements Einstein summation over arbitrary tensor networks with GPU acceleration via cuBLAS and cuTENSOR. It uses Julia's multiple dispatch to select the most efficient backend for each contraction—standard matrix multiplication for simple cases, cuTENSOR for general tensor networks—without runtime overhead. The package is especially valuable in quantum computing and machine learning research, where large tensor network contractions are a core computational primitive.

TensorOperations.jl provides fast tensor contractions, permutations, and traces using Einstein index notation, with GPU acceleration through hardware-agnostic Strided.jl implementations as well as a dedicated cuTENSOR backend. The package supports automatic differentiation and offers flexible backend selection, allowing the same high-level expression to dispatch to optimized implementations on whichever hardware is available. It is a go-to tool in quantum chemistry and condensed matter physics, where tensor operations on large arrays are ubiquitous.

Whole-program optimization

Reactant.jl takes a different approach to GPU execution: rather than offering array types or kernel abstractions, it compiles entire Julia functions to MLIR and optimizes them for execution on CPUs, GPUs, and TPUs via XLA. It uses operator tracing (aka partial evaluation) to obtain an equivalent MLIR code of the program. It then runs a ton of compiler optimizations that perform automatic differentiation, parallelization and optimization. Starting from your code written with existing packages, like CUDA.jl or KernelAbstractions.jl, Reactant will automatically perform optimizations like kernel fusion, and offload to your chosen architecture.

Reactant.jl tries to be minimally intrusive, but operator tracing may run into problems with control flow. A companion sub-package, ReactantCore.jl, exposes the @trace macro, which correctly marks control-flow constructs (if, for, etc.) during tracing. The @trace translates to a no-op if evaluated outside of the Reactant compilation context, allowing Reactant integration of the broader Julia ecosystem without fully depending on Reactant.

Task Runtimes

Dagger.jl is a Julia task runtime and scheduler that supports scalable, distributed multi-GPU execution across all 5 main GPU backends, with built-in support for KernelAbstractions-written kernels. Dagger allows the expression of generic scalable algorithms that seamlessly scale from 0 to 1 to N GPUs without having to handle the vagaries of GPU programming and state management – Dagger handles this in the background, while maximizing throughput.

AMDGPU.jl 2.6 and 2.7: linear algebra, sparse arrays, and RDNA4 matrix cores

By: Ludovic Räss

Re-posted from: https://juliagpu.org/post/2026-07-06-amdgpu-2.7/index.html

The 2.6 and 2.7 releases of AMDGPU.jl broaden the package's linear-algebra coverage — GPU Cholesky, LU and now a singular value decomposition via rocSOLVER, mixed-precision matmul, and a sparse-array interface — add WMMA matrix-core support for RDNA4 GPUs, and trim time-to-first-kernel. The documentation has also been substantially expanded.

These features span the recent 2.6 and 2.7 releases. AMDGPU.jl runs on 64-bit Linux and Windows with ROCm 6.0 or later, on Julia 1.10 through 1.13; MI300-series GPUs require Julia 1.12 or later. As always, AMDGPU.versioninfo() reports what was detected on your system.

Dense linear algebra

The dense LinearAlgebra surface on ROCArray has grown considerably. Version 2.6 added GPU Cholesky and LU factorizations (with mixed-precision support and, on Julia 1.12, the allowsingular option), and 2.7 adds Hermitian support and mixed-precision matrix multiplication.

The most recent addition is a singular value decomposition: svd, svd!, svdvals and cond now work directly on ROCMatrix (#960), backed by rocSOLVER. Previously these fell through to LinearAlgebra's default divide-and-conquer path, which rocSOLVER does not implement. Two algorithms are available through an alg keyword — a QR iteration (QRAlgorithm, gesvd!) and a one-sided Jacobi method (JacobiAlgorithm, gesvdj!) — with Jacobi the default, as it is consistently faster on AMD hardware.

julia> using AMDGPU, LinearAlgebrajulia> A = AMDGPU.rand(Float32, 1000, 1000);julia> F = svd(A);          # Jacobi by default; also svdvals, condjulia> L = cholesky(A'A);   # and lu, qr, \, mul!

The two algorithms differ noticeably in practice. The table below shows the time for a full SVD of an n×n Float32 matrix on an MI300X, measured by Evelyne Ringoot (@evelyne-ringoot) (times in milliseconds; see #837 for the full data):

size (n×n) QR iteration Jacobi
256 165 40
1024 2,128 273
4096 30,508 2,989
8192 122,498 16,555

Sparse arrays

Sparse arrays gained a linear-algebra interface on top of rocSPARSE. The ROCSparseMatrixCSR, ROCSparseMatrixCSC and ROCSparseMatrixCOO types convert to and from a host SparseMatrixCSC, and sparse matrix–vector and matrix–matrix products work through the standard * operator, alongside format conversions and preconditioner building blocks.

FFTs

The rocFFT plan and execution path was redesigned for more predictable handling of real and complex transforms and of in-place versus out-of-place plans, all through the standard AbstractFFTs.jl interface (fft, plan_fft, rfft, and friends).

Matrix cores on RDNA4

Version 2.6 adds WMMA (Wave Matrix Multiply-Accumulate) support for RDNA4 / gfx1201+ GPUs (#929, by @ffrancesco94), alongside the existing RDNA3 support. The intrinsics live in the AMDGPU.Device.WMMA_RDNA3 and AMDGPU.Device.WMMA_RDNA4 submodules; for the 2.x cycle, WMMA aliases the RDNA3 module (#955).

A leaner load and broader toolchain support

A precompilation workload was added to cut time-to-first-kernel, and SpecialFunctions was moved into a package extension so it is only loaded when actually used. On the toolchain side, recent releases track the Julia 1.13 device libraries and LLVM up to 21.1, and make device discovery more robust on different Linux distributions. Much of this compiler and runtime work builds on the shared GPU infrastructure maintained by Valentin Churavy (@vchuravy), Gabriel Baraldi (@gbaraldi) and others.

Documentation

The documentation has been substantially expanded (#959): new usage guides for array programming, tasks and streams, and KernelAbstractions; a Libraries section covering rocBLAS/rocSOLVER, rocSPARSE, rocFFT, rocRAND and MIOpen; an FAQ with guidance on depending on AMDGPU.jl conditionally; and a feature overview on the documentation home page.

As always, update to the latest version to pick these up, and see the changelog for the full list. AMDGPU.jl is largely a community effort, and contributions, issue reports and feedback are all welcome. Thanks to Evelyne Ringoot, @ffrancesco94, Valentin Churavy, Gabriel Baraldi, and everyone else who contributed to these releases.

Metal.jl 1.10: Linear algebra, FFTs, and a faster runtime

By: Christian Guinard, Tim Besard

Re-posted from: https://juliagpu.org/post/2026-07-01-metal-1.10/index.html

Metal.jl 1.10 is a big release. It adds native matrix multiplication, GPU-accelerated linear solvers and FFTs, BFloat16 support, and MPS-backed reductions, scans and sorting. The runtime also got considerably faster and leaner, and there is a new in-process profiler.

Before getting into the new features, one thing to flag up front: Metal.jl 1.10 requires macOS 14 or later, up from macOS 13. On older systems the package now refuses to initialize, and Metal.functional() returns false. The supported range is macOS 14 through 26, on Julia 1.10 through 1.13.

Tied to that requirement is a change in how kernels are compiled. Previously Metal.jl pinned a conservative baseline (AIR 2.5 / metallib v1.2.6) regardless of the host. Since Metal.jl only ever compiles for the machine it runs on, it now emits the newest AIR, MSL and metallib versions the host macOS supports, exactly like Apple's offline metal compiler does. That unlocks newer language features for free: AIR 2.6 / Metal 3.1 on macOS 14, up to AIR 2.8 / Metal 4.0 on macOS 26 (and Metal 4.1 on the macOS 27 beta). You can see what your machine targets in versioninfo:

julia> Metal.versioninfo()
macOS 26.6.0, Darwin 25.6.0Toolchain:
- Julia: 1.12.6
- LLVM: 18.1.7
- Metal: 4.0 (MSL), 2.8 (AIR), 1.2.9 (metallib)Julia packages:
- Metal.jl: 1.10.0
- GPUArrays: 11.5.8
- GPUCompiler: 1.22.7
- KernelAbstractions: 0.9.42
- ObjectiveC: 6.0.0
- LLVM: 9.10.0
- LLVMDowngrader_jll: 0.8.1+01 device:
- Apple M3 Pro (14 GPU cores, 80.000 KiB allocated; Apple9, Metal4 family)

Native matrix multiplication

Up to now, every A * B on an MtlArray went straight to Apple's vendor libraries (Metal Performance Shaders or MPSGraph). That works well on large matrices, but it leaves us at the mercy of the vendor: there are eltypes MPS does not support, small matrices pay a steep launch overhead, and bugs like the M1/M2 matmul NaN issue are out of our hands. Anything unsupported fell back to GPUArrays' generic implementation, resulting in poor performance.

Metal.jl 1.10 ships its own native GEMM kernels. You pick a backend through the Metal.matmul_alg scoped value, which defaults to :auto:

  • :scalar is a per-element tiled kernel that handles any Metal-supported eltype (integers, complex, BFloat16) and any transpose or offset. It's the universal fallback.

  • :simd is a simdgroup_matrix kernel for Float16/Float32 (and BFloat16) with Float32 accumulation

  • :tensor is a Metal 4 tensor_ops::matmul2d kernel, available on Metal 4-capable devices running macOS 26+.

  • :native picks the best of the three, per device and per operand.

  • :auto (the default) tries the vendor libraries first, then falls back to :native.

Linear solvers

Closing a long-standing request, many more standard LinearAlgebra operations on Float32/Float16 MtlMatrixes now run on the GPU through MPS-backed solvers. That covers \, lu, cholesky (including on Symmetric/Hermitian wrappers), triangular solves, and inv/det/logdet:

julia> using Metal, LinearAlgebrajulia> A = MtlArray(rand(Float32, 512, 512) + 512I);julia> b = MtlArray(rand(Float32, 512));julia> x = A \ b;                              # MPS LU solve, on the GPUjulia> norm(Array(A) * Array(x) - Array(b))    # residual, at the Float32 noise floor
4.4064095f-6julia> M = MtlArray(rand(Float32, 256, 256));julia> logdet(cholesky(Symmetric(M'M + I)))    # cholesky factorization, also on the GPU
638.44586f0

BFloat16

BFloat16 arrays now run natively on the GPU as well:

julia> using Metal, BFloat16sjulia> a = MtlArray(BFloat16[1.5, 2.5, 3.5])
3-element MtlVector{BFloat16, Metal.PrivateStorage}:
 1.5
 2.5
 3.5julia> sum(a .* BFloat16(2))
BFloat16(15.0f0)

All Julia versions are supported, but before Julia 1.13 operations involving scalar BFloat16 values (e.g. a .+ BFloat16(1)) may be slower because they go through a software emulation path in BFloats.jl.

FlashAttention example

To tie the new building blocks together, there is a FlashAttention example that spells out scaled dot-product attention in four different ways, one per programming model Metal.jl exposes:

  • with plain array operations (*, broadcasting, maximum, sum, exp);

  • with MPSGraph's fused scaledDotProductAttention op;

  • with a hand-written kernel built on MtlSimdgroupMatrix{Float16,8,8};

  • and with a fused kernel using the Metal 4 tensor_ops::matmul2d primitives.

It's a good read if you want to see how the simdgroup and tensor intrinsics look in practice; you'll find it in examples/flashattention.jl.

Reductions, scans and sorting

Reductions, prefix scans and sorting now route through MPSGraph when it makes sense. This speeds up reductions and scans, and introduces support for sorting:

julia> sort(MtlVector(Int16[5, -3, 2, 9, -7, 0]))
6-element MtlVector{Int16, Metal.PrivateStorage}:
 -7
 -3
  0
  2
  5
  9julia> accumulate(max, MtlVector(Int32[1, 3, 2, 5, 4]))
5-element MtlVector{Int32, Metal.PrivateStorage}:
 1
 3
 3
 5
 5

Neural-network primitives

Metal.jl 1.10 also wraps the core MPSGraph neural-network primitives: softmax and logsoftmax, 2D convolution, and max/mean pooling, each with its gradient. These are wired up as the Metal backend for NNlib.jl, so once that release lands, Flux models gain GPU acceleration on Apple hardware through the functions you already use (conv, maxpool, softmax, …) rather than any Metal-specific API.

FFTs

On the back of the MPSGraph work, Metal.jl now supports FFTs through the AbstractFFTs.jl interface:

julia> using Metal, AbstractFFTsjulia> x = MtlArray(rand(ComplexF32, 2048, 2048));julia> y = fft(x);          # just worksjulia> Array(ifft(y)) ≈ Array(x)
truejulia> p = plan_fft(x);     # reusable plans, toojulia> Array(p * x) ≈ Array(y)
true

Real transforms (rfft/irfft), transforms along specific dimensions, and batched transforms are all supported. Running on the GPU is a large win over a CPU FFT, even one backed by AppleAccelerate. The following are timings on a 30-core M2 Max:

Size CPU (FFTW) CPU (FFTW + AppleAccelerate) GPU (Metal) speedup vs. Accelerate
512×512 4.2 ms 766.2 µs 173.4 µs 4.4×
1024×1024 19.7 ms 3.7 ms 246.3 µs 15×
2048×2048 99.5 ms 20.8 ms 588.4 µs 35×
4096×4096 580.1 ms 99.0 ms 2.5 ms 39×

A faster, leaner runtime

A lot of work in this cycle went into the cost of getting work onto the GPU and back.

Batched command submission. Metal.jl used to create, encode and commit a fresh command buffer for every single launch. It now keeps one command buffer open and submits launches into it, flushing on synchronization or other triggers. That amortizes the per-launch command-buffer overhead, which is the dominant cost for workloads built out of many small kernels.

Non-blocking synchronization. Synchronization was ported from CUDA.jl to a spin-then-yield scheme instead of blocking inside Metal. The primary motivation is correctness (a blocked main thread can deadlock against a Metal callback that needs to do I/O), but it is also dramatically faster on the fast paths:

Scenario before after speedup
synchronize() on a queue that never ran work 15.87 µs 0.19 µs ~86×
synchronize() when the queue is idle 15.55 µs 0.37 µs ~42×
small kernel + synchronize() in a tight loop 359 µs 149 µs ~2.4×

GC under memory pressure. Because MtlArray buffers are allocated by Metal, Julia's garbage collector can't see them, and on a unified-memory Mac that means it happily lets you allocate until the system starts paging and freezes. Metal.jl now reads the memory pressure straight from Metal and triggers an incremental GC when usage gets high (above 75% normally, lower on synchronization points where the pause is hidden behind a wait anyway), rate-limited so it never spends more than a small fraction of wall-clock time collecting.

Cheaper object lifetimes. The hand-rolled retain/release/finalizer bookkeeping for Metal objects was replaced with ObjectiveC.jl's automatic reference counting, removing a few hundred lines of fiddly code and simplifying per-launch bookkeeping .

Faster large copies. Shared-storage GPU→GPU copies used to always go through a CPU memcpy. For large arrays it's faster to use a GPU blit, so copies above 32 MB now switch to that path (small copies stay on memcpy, where the API overhead would dominate):

Size before (CPU memcpy) after speedup
64 MB 3.28 ms 1.19 ms 2.8×
256 MB 6.54 ms 2.08 ms 3.1×
1024 MB 21.55 ms 5.98 ms 3.6×

Separately, copies larger than 4 GiB no longer silently fail; they are chunked into pieces Metal can handle.

Time to first kernel. A real precompilation workload plus some despecialization brought the time to a first kernel down significantly:

$ julia -e 'using Metal; a = MtlArray([1, 2, 3]); @time a .+ 1'
0.161035 seconds (178.09 k allocations: 8.552 MiB, 52.46% compilation time: 22% of which was recompilation)

Compare that to the previous version of Metal.jl:

$ julia -e 'using Metal; a = MtlArray([1, 2, 3]); @time a .+ 1'
8.133787 seconds (33.91 M allocations: 1.636 GiB, 3.64% gc time, 99.25% compilation time: 1% of which was recompilation)

A profiler that doesn't need Xcode

Timing a single kernel with BenchmarkTools is easy enough, but understanding where time goes in a larger program used to mean reaching for Xcode's Instruments. Metal.jl 1.10 adds an in-process profiler, Metal.@profile, that captures the GPU operations Metal.jl submits and prints a summary, no Xcode required:

julia> a = Metal.rand(Float32, 1024, 1024); b = similar(a); c = similar(a);julia> b .= a .+ 1f0; c .= sqrt.(b); Metal.synchronize();   # warm upjulia> Metal.@profile begin
           b .= a .+ 1f0
           c .= sqrt.(b)
       end
Profiled over 58.7 ms.Host-side activity: 42 Objective-C calls taking 118.0 µs (0.20% of wall-clock)
┌──────────┬────────────┬───────┬──────────────────────────────────────────┐
│ Time (%) │ Total time │ Calls │ Name                                     │
├──────────┼────────────┼───────┼──────────────────────────────────────────┤
│    0.06% │   33.29 µs │     2 │ [MTLCommandBuffer commit]                │
│    0.04% │   22.17 µs │     2 │ [MTLCommandQueue commandBuffer]          │
│    0.04% │    22.0 µs │     2 │ [MTLCommandBuffer computeCommandEncoder] │
│     ...  │     ...    │   ... │ ...                                      │
└──────────┴────────────┴───────┴──────────────────────────────────────────┘Device-side activity: GPU was busy 831.75 µs (1.42% of wall-clock)
┌──────────┬────────────┬───────┬───────────────────────────┬──────────────┐
│ Time (%) │ Total time │ Calls │ Time distribution         │ Name         │
├──────────┼────────────┼───────┼───────────────────────────┼──────────────┤
│    1.42% │  831.75 µs │     2 │ 415.87 µs ± 220.62        │ broadcast_2d │
└──────────┴────────────┴───────┴───────────────────────────┴──────────────┘

The host table groups the Objective-C calls, the device table groups kernels and blits, and the slowest operations are color-highlighted. Pass trace=true for a chronological timeline (with threadgroup, occupancy and threadgroup-memory columns) instead of a summary, and use Metal.@bprofile to benchmark a snippet by running it repeatedly. The old Xcode-based capture is still there under Metal.@profile external=true. One caveat worth knowing: MPS and MPSGraph operations, including the default matmul backend, submit their own command buffers and don't show up in the integrated trace yet, so reach for the external profiler to inspect those.

Better debugging

Device-side printing. On macOS 15+, you can now print from inside a kernel, built on Apple's os_log. There's @mtlprintf, plus the friendlier @mtlprint, @mtlprintln, and @mtlshow:

julia> function device_println()
           @mtlprintln("Hello, world!")
           return
       endjulia> @metal device_println();
Hello, world!

It also wires up KernelAbstractions' @print, so the same works in KA kernels.

Richer exceptions. When a kernel throws, say a bounds error, that used to surface as an opaque failure. Device exceptions are now reported back to the host as a KernelException carrying the actual cause:

julia> function oob(a)
           a[2] = 1f0   # a has length 1
           return
       endjulia> a = MtlArray(zeros(Float32, 1));julia> @metal threads=1 oob(a)
ERROR: KernelException: A BoundsError was thrown

Launching with debug_level=2 adds a full device-side stacktrace. The detailed machinery only kicks in at the higher debug level, so the common case stays fast.

Device-side allocation. A minimal device-side malloc means kernels that need dynamic allocation (notably exception-throwing code and some broadcasts) now compile and run where they previously failed outright.

New intrinsics

The warp-level primitive set is now much more complete. Metal.jl 1.10 adds the indexed simd_shuffle/simd_shuffle_xor shuffles, the simd_ballot/simd_vote_all/simd_vote_any voting intrinsics, and the full set of quad-group (4-thread) equivalents: quad_shuffle, quad_ballot, quad_vote_all, and friends. There are also UInt16 variants of every thread- and grid-indexing intrinsic (thread_position_in_grid_i16() and so on) for when 16-bit indices are enough.

Please refer to the Metal Shading Language Specification to verify where indexing types must match for a kernel to be valid.

Other improvements

Metal.jl 1.10 includes plenty more:

One breaking fix to be aware of: launching a kernel with a grid dimension larger than typemax(UInt32) used to silently truncate. It now raises an error instead, so use grid-stride loops for kernels that need to cover arrays larger than that.

As always, update to the latest version to get these improvements, and check out the changelog for the full list.