Compiler

@metal [kwargs...] f(args...)

Compile f(args...) as a GPU kernel and launch it on the current device. f must return nothing. The kernel is compiled on first use for the types of f and args, and cached; arguments are converted to their device counterparts with mtlconvert (e.g., an MtlArray becomes an MtlDeviceArray).

The launch is asynchronous: it is recorded on the command queue of the current task, and submitted together with other work. Use synchronize to wait for it, which is also where errors thrown by the kernel are reported (see Metal.KernelException).

@metal dispatches groups threadgroups of threads threads each, so the grid holds groups .* threads threads per dimension. If that is more than the number of elements to process, the kernel has to check bounds itself, e.g., i <= length(a) || return.

Launch keywords

• threads (default 1): the number of threads per threadgroup, an integer or a tuple of up to three integers. The product cannot exceed the kernel's limit, which is at most 1024 on Apple GPUs but can be lower for complex kernels: compile the kernel with launch=false and check its maxthreads (see HostKernel).
• groups (default 1): the number of threadgroups in the grid, an integer or a tuple of up to three integers.
• queue: the command queue to submit to, instead of the task's global_queue.
• submit (default false): submit the queue's pending work right after this launch, instead of when the batch is full or the queue is synchronized.
• launch (default true): when false, only compile the kernel and return it as a HostKernel, to be launched later by calling it with arguments and launch keywords. This is not allowed together with other launch keywords.

Compiler keywords

• name: the name of the kernel in the generated code, e.g., as shown in Xcode. Defaults to the name of f.
• fastmath (default false, or true when Julia runs with --math-mode=fast): apply @fastmath semantics to all floating-point code in the kernel, and compile it with Metal's fast-math mode.
• debug_level (default: Julia's -g level): how much information is kept to report device-side exceptions; 2 adds the position of the faulting thread and a backtrace.
• opt_level (default 2): the optimization level, from 0 to 3, independent of the host session's -O level.
• always_inline (default false): inline every function called by the kernel.
• minthreads: the exact threadgroup size the kernel will be launched with, an integer or a tuple of up to three integers, like MSL's required_threads_per_threadgroup. The compiler can specialize for it, and launches with any other size are rejected. Requires macOS 26.
• maxthreads: the maximum number of threads per threadgroup the kernel will be launched with, like MSL's max_total_threads_per_threadgroup. A lower maximum leaves the compiler more registers per thread, e.g., to avoid running out of them in complex kernels.
• macos, metal, air, gpufamily: the deployment target, Metal Shading Language version, AIR version and Apple GPU family to compile for, which default to what the host supports. See mtlfunction.

Examples

function vadd(c, a, b)    i = thread_position_in_grid().x    if i <= length(c)        @inbounds c[i] = a[i] + b[i]    end    returnenda = Metal.rand(Float32, 10_000); b = Metal.rand(Float32, 10_000)c = similar(a)threads = 256@metal threads groups=cld(length(c), threads) vadd(c, a, b)

HostKernel

A compiled kernel, as returned by mtlfunction or @metal launch=false. Launch it by calling it (see below). Its fields describe the compute pipeline Metal created for it:

• maxthreads: the maximum number of threads per threadgroup this kernel can be launched with (maxTotalThreadsPerThreadgroup), which can be lower than the device's limit of 1024 for kernels that use many registers.
• exec_width: the thread execution width, i.e., the number of threads in a SIMD-group (threadExecutionWidth). Threadgroup sizes that are a multiple of it use the GPU best.
• tgmem: the amount of threadgroup memory, in bytes, that the kernel allocates statically (with MtlThreadGroupArray).
• pipeline: the MTLComputePipelineState, for other properties.

kernel = @metal launch=false vadd(c, a, b)threads = min(length(c), kernel.maxthreads)kernel(c, a, b; threads, groups=cld(length(c), threads))

mtlfunction(f, tt=Tuple{}; kwargs...)::HostKernel

Compile f for argument types tt (the types after mtlconvert) for the current device, and return the kernel as a HostKernel. This is what @metal uses; call it directly to compile without converting arguments.

The compiler keywords of @metal are supported, and:

• macos and metal: the deployment target (macOS 15 or later) and Metal Shading Language version to compile for, which default to the newest the host supports (Metal.metal_target). A language version that needs a newer macOS raises the deployment target. Code for versions the host doesn't support can be generated and inspected, but not loaded. Targeting Metal 3.1 or older rules out barrier and fence scopes, and ordered atomics.
• air: the AIR bitcode version, which otherwise follows the deployment target. Meant for testing.
• gpufamily: the Apple GPU family (MTL.MTLGPUFamilyApple<n>) that the generated code may rely on (see apple_family). Defaults to the highest family the device supports.
• source: the callable before mtlconvert turned it into f, which is converted again at every launch (so it should convert to an object of the same type as f). Defaults to f, but should be set when f captures GPU arrays, so that their buffers are made available to the GPU and kept alive as long as the kernel is in use.

The result is cached, so calling mtlfunction in a hot path is cheap. Code is regenerated when f is redefined, or for different types or keyword arguments.

mtlconvert(x)

Convert a kernel argument x to the object the kernel receives, e.g., an MtlArray to an MtlDeviceArray. @metal calls it on every argument. Tuples, named tuples and array wrappers are converted recursively; other objects are passed as-is.

To convert the fields of your own types, don't add methods to this function, but use Adapt.jl (e.g., Adapt.@adapt_structure MyType), which all JuliaGPU back-ends share.

metalcall(pipeline::MTLComputePipelineState, types, values...;
          groups=1, threads=1, queue=nothing, submit=false,
          threadgroup_memory=(), resources=())

ccall-like interface for launching a compute pipeline, e.g., one created from a function in a precompiled Metal library, on the GPU. types is a tuple of types, or a tuple type, with an entry per parameter of the kernel, which determines how each of the values is passed:

• pointer types (Metal.MtlPtr, Core.LLVMPtr or Ptr) are bound as a buffer (with an offset), and should be passed an MtlArray, MtlPtr or MTLBuffer;
• other types are passed by value, after converting the argument with mtlconvert and convert; these must be bits types with the layout the kernel expects.

The parameters are bound to consecutive buffer indices, starting at 0, skipping those of a type that holds no data: a singleton type like Nothing, or a type like Type{Int}. Nothing is checked against the kernel's signature: using types that don't match it is undefined behavior.

The kernel is launched on groups threadgroups of threads threads each, both of which should be an integer or a tuple of up to three integers. The other keyword arguments are:

• queue: the command queue to launch on, defaulting to the current task's;
• submit: submit the work batched on the queue right away, instead of when Metal.jl decides to;
• threadgroup_memory: the length in bytes of the threadgroup memory for each of the kernel's [[threadgroup(i)]] parameters, in order (a multiple of 16 bytes);
• resources: buffers that the kernel accesses through GPU addresses it reads from its arguments, e.g., from a struct or an argument buffer. These are made resident and kept alive until the kernel has finished, which happens automatically for the buffers of arrays in arguments converted by mtlconvert.

For example, to launch a kernel written in the Metal Shading Language:

lib = MTL.MTLLibrary(device(), """    kernel void scale(device float *x [[buffer(0)]], constant float &a [[buffer(1)]],                      uint i [[thread_position_in_grid]]) {        x[i] *= a;    }    """)pipeline = MTL.MTLComputePipelineState(device(), MTL.MTLFunction(lib, "scale"))x = Metal.ones(Float32, 256)metalcall(pipeline, Tuple{Metal.MtlPtr{Float32}, Float32}, x, 2; threads=256)

See also Metal.metalcall_tuple, which takes the values as a tuple.

metalcall_tuple(pipeline, types, values::Tuple; kwargs...)

The same as metalcall, but taking the values as a tuple, which is faster for kernels with many arguments.

@device_code_lowered ex

Evaluate ex, and return the lowered Julia IR of every kernel it compiles, like InteractiveUtils.@code_lowered.

@device_code_typed [kwargs...] ex

Evaluate ex, and return the typed Julia IR of every kernel it compiles, like InteractiveUtils.@code_typed. To debug a kernel that fails to compile, pass only_failed=true interactive=true, which inspects that kernel with Cthulhu.jl.

@device_code_warntype [io::IO=stdout] [kwargs...] ex

Evaluate ex, and print the typed Julia IR of every kernel it compiles, highlighting values whose type could not be inferred, like InteractiveUtils.@code_warntype.

@device_code_llvm [io::IO=stdout] [kwargs...] ex

Evaluate ex, and print the LLVM IR that Julia generates for every kernel it compiles, before it is lowered to AIR (see @device_code_air). See Metal.code_llvm for the keyword arguments.

@device_code_air [io::IO=stdout] [kwargs...] ex

Evaluate ex, and print the AIR of every kernel it compiles: the LLVM IR after lowering to AIR intrinsics and downgrading to the bitcode format Metal reads, disassembled again. This is the code that Metal compiles for the GPU.

See Metal.code_air for the keyword arguments.

@device_code_metal [io::IO=stdout] [kwargs...] ex

An alias for @device_code_air.

@device_code dir::AbstractString=... ex

Evaluate ex, and write all intermediate forms of the code of every kernel it compiles to files in the directory dir.

Metal.code_typed(f, types; kernel=false, kwargs...)

Like InteractiveUtils.code_typed, return the typed Julia IR for calling f with arguments of types on the GPU. See Metal.code_air for the arguments.

Metal.code_warntype([io=stdout], f, types; kernel=false, kwargs...)

Like InteractiveUtils.code_warntype, print the typed Julia IR for calling f with arguments of types on the GPU, highlighting values whose type could not be inferred. See Metal.code_air for the arguments.

Metal.code_llvm([io=stdout], f, types; kernel=false, kwargs...)

Like InteractiveUtils.code_llvm, print the LLVM IR that Julia generates for calling f with arguments of types on the GPU, before it is lowered to AIR. See Metal.code_air for the arguments.

Metal.code_air([io=stdout], f, types; kernel=false, kwargs...)

Print the AIR that Metal.jl generates for calling f with arguments of types on the GPU: the LLVM IR after lowering to AIR intrinsics and downgrading to the bitcode format Metal reads, disassembled again. This is the code that Metal compiles for the GPU.

The types are those of the arguments inside the kernel, e.g., MtlDeviceVector{Float32,1} rather than MtlVector{Float32}, so the @device_code_air macro is usually easier to use. Pass kernel=true to compile f as a kernel, like @metal does, and compiler keywords of @metal such as metal to change the target. Other keywords are passed to GPUCompiler.jl.

Metal.code_native([io=stdout], f, types; kwargs...)

An alias for Metal.code_air: AIR is the closest to native code that Metal.jl generates.