GitHub

Kernel functions

Functions available inside kernels, grouped by topic. See Kernel programming and Advanced kernel programming for examples, and the Metal Shading Language specification for the corresponding Metal operations. Julia-specific behavior is documented below.

Indexing

The positions and sizes of the thread, its SIMD-group, quad-group and threadgroup, and the grid. Positions and indices are 1-based. See Apple's Creating threads and threadgroups.

Metal.thread_position_in_grid — Function
thread_position_in_grid()::@NamedTuple{x::UInt32, y::UInt32, z::UInt32}

Return the 1-based position of the current thread in the grid, per dimension. For a one-dimensional launch, thread_position_in_grid().x is the global index of the thread.

See also thread_position_in_grid_i16, which returns UInt16 values.

source
Metal.thread_position_in_threadgroup — Function
thread_position_in_threadgroup()::@NamedTuple{x::UInt32, y::UInt32, z::UInt32}

Return the 1-based position of the current thread in its threadgroup, per dimension.

See also thread_position_in_threadgroup_i16, which returns UInt16 values.

source
Metal.threadgroup_position_in_grid — Function
threadgroup_position_in_grid()::@NamedTuple{x::UInt32, y::UInt32, z::UInt32}

Return the 1-based position of the current threadgroup in the grid, per dimension.

See also threadgroup_position_in_grid_i16, which returns UInt16 values.

source
Metal.threads_per_grid — Function
threads_per_grid()::@NamedTuple{x::UInt32, y::UInt32, z::UInt32}

Return the size of the grid in threads, per dimension: the product of the groups and threads arguments of @metal.

See also threads_per_grid_i16, which returns UInt16 values.

source
Metal.threads_per_threadgroup — Function
threads_per_threadgroup()::@NamedTuple{x::UInt32, y::UInt32, z::UInt32}

Return the number of threads in the current threadgroup, per dimension: the threads argument of @metal.

See also threads_per_threadgroup_i16, which returns UInt16 values.

source
Metal.threadgroups_per_grid — Function
threadgroups_per_grid()::@NamedTuple{x::UInt32, y::UInt32, z::UInt32}

Return the number of threadgroups in the grid, per dimension: the groups argument of @metal.

See also threadgroups_per_grid_i16, which returns UInt16 values.

source
Metal.thread_index_in_threadgroup — Function
thread_index_in_threadgroup()::UInt32

Return the 1-based linear index of the current thread in its threadgroup.

See also thread_index_in_threadgroup_i16, which returns UInt16 values.

source
Metal.thread_index_in_simdgroup — Function
thread_index_in_simdgroup()::UInt32

Return the 1-based index of the current thread in its SIMD-group (its lane).

See also thread_index_in_simdgroup_i16, which returns UInt16 values.

source
Metal.thread_index_in_quadgroup — Function
thread_index_in_quadgroup()::UInt32

Return the 1-based index of the current thread in its quad-group (1 to 4).

See also thread_index_in_quadgroup_i16, which returns UInt16 values.

source
Metal.simdgroup_index_in_threadgroup — Function
simdgroup_index_in_threadgroup()::UInt32

Return the 1-based index of the current SIMD-group in its threadgroup.

See also simdgroup_index_in_threadgroup_i16, which returns UInt16 values.

source
Metal.quadgroup_index_in_threadgroup — Function
quadgroup_index_in_threadgroup()::UInt32

Return the 1-based index of the current quad-group in its threadgroup.

See also quadgroup_index_in_threadgroup_i16, which returns UInt16 values.

source
Metal.simdgroups_per_threadgroup — Function
simdgroups_per_threadgroup()::UInt32

Return the number of SIMD-groups in the current threadgroup.

See also simdgroups_per_threadgroup_i16, which returns UInt16 values.

source
Metal.quadgroups_per_threadgroup — Function
quadgroups_per_threadgroup()::UInt32

Return the number of quad-groups in the current threadgroup.

See also quadgroups_per_threadgroup_i16, which returns UInt16 values.

source
Metal.threads_per_simdgroup — Function
threads_per_simdgroup()::UInt32

Return the number of threads in a SIMD-group, i.e., the thread execution width (32 on current Apple GPUs).

See also threads_per_simdgroup_i16, which returns UInt16 values.

source
Metal.thread_execution_width — Function
thread_execution_width()::UInt32

Return the thread execution width, i.e., the number of threads in a SIMD-group. Deprecated in MSL; use threads_per_simdgroup instead.

See also thread_execution_width_i16, which returns UInt16 values.

source
Metal.grid_origin — Function
grid_origin()::@NamedTuple{x::UInt32, y::UInt32, z::UInt32}

Return the 1-based origin of the grid, for kernels that read per-thread stage-in data.

See also grid_origin_i16, which returns UInt16 values.

source
Metal.grid_size — Function
grid_size()::@NamedTuple{x::UInt32, y::UInt32, z::UInt32}

Return the size of the grid, for kernels that read per-thread stage-in data.

See also grid_size_i16, which returns UInt16 values.

source

16-bit variants

Versions and GPU families

Metal.metal_version — Function
metal_version()

Return the Metal Shading Language version that the kernel is compiled for, as a version that compares with sv"..." literals. This is a compile-time constant, so a branch on it is resolved during compilation and can guard code that needs a newer version:

if metal_version() >= sv"4.1"    # use functionality that needs MSL 4.1end

The version defaults to the newest the host supports (see Metal.metal_target), and can be chosen with the metal keyword of @metal. Only use this in kernels.

source
Metal.air_version — Function
air_version()

Return the AIR version that the kernel is compiled for, as a version that compares with sv"..." literals. Like metal_version, this is a compile-time constant. Only use this in kernels.

source
Metal.apple_family — Function
apple_family()

Return the n of the highest MTLGPUFamilyApple<n> that the kernel is compiled for (e.g., 7 for M1, 8 for M2, 9 for M3 and M4, 10 for M5), or 0 if the device reports no Apple family. Like metal_version, this is a compile-time constant that can guard family-specific code. Only use this in kernels.

source

Device arrays

Metal.MtlDeviceArray — Type
MtlDeviceArray{T,N,A} <: DenseArray{T,N}

The array type that kernels see: an N-dimensional, column-major array of elements of type T in address space A (device memory for the MtlArrays passed to a kernel, threadgroup memory for MtlThreadGroupArrays). @metal converts MtlArray arguments to it. It supports indexing, iteration and contiguous views, and is only valid inside the kernel.

MtlDeviceArray(dims, ptr)MtlDeviceArray{T}(dims, ptr)MtlDeviceArray{T,N}(dims, ptr)MtlDeviceArray{T,N,A}(dims, ptr, [maxsize])

Wrap the device pointer ptr::Core.LLVMPtr{T,A} in an array of size dims, an integer or a tuple. T and A follow from the type of ptr, and N from the length of dims; when given, they have to match.

Arrays of isbits unions (e.g. Union{Int32,Nothing}) keep one selector byte per element past their data, like Array: maxsize is the distance in bytes from ptr to those bytes, which defaults to the size of the data.

source
Metal.Const — Type
Metal.Const(A::MtlDeviceArray)

Mark a device array as read-only for the duration of the kernel, which lets the compiler optimize loads from it (e.g., hoist them out of loops). Base.Experimental.Const(A) does the same. This interface is experimental, and might change without deprecation.

Warning

Writing to the underlying memory while the kernel executes, through A or another array, is undefined behavior, and can produce wrong results.

source

Threadgroup memory

Metal.MtlThreadGroupArray — Function
MtlThreadGroupArray(T, dims)

Allocate an array of element type T and size dims in threadgroup memory, and return it as an MtlDeviceArray. Every threadgroup gets its own array, shared by its threads. The size has to be a compile-time constant (e.g., a literal, or derived from type parameters), the memory is uninitialized, and a kernel can use at most 32 KiB of threadgroup memory.

Use threadgroup_barrier with Metal.MemoryFlagThreadGroup between writing the array and reading what other threads wrote:

function reverse_kernel(a)    i = thread_position_in_threadgroup().x    tmp = MtlThreadGroupArray(Float32, 64)    tmp[i] = a[i]    threadgroup_barrier(Metal.MemoryFlagThreadGroup)    a[i] = tmp[65 - i]    returnend
source

Synchronization

Barriers and fences, which order memory accesses between threads. See Apple's Resource synchronization.

Metal.threadgroup_barrier — Function
threadgroup_barrier(flags=Metal.MemoryFlagNone, order=Metal.memory_order_seq_cst,
                    scope=Metal.thread_scope_threadgroup)
threadgroup_barrier(flags, scope)

Wait until all threads in the threadgroup reach the barrier, and order their accesses to the memory selected by flags (see MemoryFlags) with memory order order (see memory_order) for the threads in scope (see thread_scope). Without flags, the barrier only synchronizes execution.

All threads in the threadgroup have to reach the same barrier, so don't call it from a branch that only some threads take. The arguments have to be compile-time constants.

source
Metal.simdgroup_barrier — Function
simdgroup_barrier(flags=Metal.MemoryFlagNone, order=Metal.memory_order_seq_cst,
                  scope=Metal.thread_scope_simdgroup)
simdgroup_barrier(flags, scope)

Wait until all threads in the SIMD-group reach the barrier, and order their accesses to the memory selected by flags (see MemoryFlags) with memory order order (see memory_order) for the threads in scope (see thread_scope). Without flags, the barrier only synchronizes execution. The arguments have to be compile-time constants.

source
Metal.atomic_thread_fence — Function
atomic_thread_fence(flags, order, scope=Metal.thread_scope_device)

Order the current thread's accesses to the memory selected by flags (see MemoryFlags) with memory order order (see memory_order), with respect to the threads in scope (see thread_scope). Unlike a barrier, a fence does not wait for other threads. The arguments have to be compile-time constants.

source
Metal.MemoryFlags — Type
MemoryFlags

Flags that select the memory a barrier or fence orders accesses to, combined with |:

  • Metal.MemoryFlagNone: no memory; a barrier only synchronizes execution.
  • Metal.MemoryFlagDevice: device memory, e.g., MtlArrays passed to the kernel.
  • Metal.MemoryFlagThreadGroup: threadgroup memory (MtlThreadGroupArray).
  • Metal.MemoryFlagTexture: texture memory, for textures with read-write access.
  • Metal.MemoryFlagThreadGroup_ImgBlock: threadgroup imageblock memory.

The barriers default to MemoryFlagNone, so to make writes to threadgroup memory visible to the other threads of the threadgroup, pass Metal.MemoryFlagThreadGroup:

threadgroup_barrier(Metal.MemoryFlagThreadGroup)

The flags have to be compile-time constants. See mem_flags in the Metal Shading Language specification.

source
Metal.memory_order — Type
memory_order

The memory order of an atomic operation, fence or barrier, as in MSL and C++:

  • Metal.memory_order_relaxed: no ordering of other memory accesses.
  • Metal.memory_order_acquire: later accesses cannot move before this operation.
  • Metal.memory_order_release: earlier accesses cannot move after this operation.
  • Metal.memory_order_acq_rel: both acquire and release.
  • Metal.memory_order_seq_cst: acquire and release, and a single total order of all sequentially consistent operations.

The order has to be a compile-time constant. Before Metal 4.1 (macOS 27), barriers and ordered fences are always sequentially consistent.

source
Metal.thread_scope — Type
thread_scope

The set of threads that a fence or barrier synchronizes with, as in MSL: Metal.thread_scope_thread, Metal.thread_scope_simdgroup, Metal.thread_scope_threadgroup or Metal.thread_scope_device. The scope has to be a compile-time constant.

source

SIMD-group functions

Functions that exchange data between the threads of a SIMD-group.

Metal.simd_shuffle — Function
simd_shuffle(data::T, lane::Integer)

Return data of the thread in the SIMD-group whose 1-based index is lane (see thread_index_in_simdgroup). Different threads can read from different lanes. T can be Float32, Float16, Int32, UInt32, Int16, UInt16, Int8 or UInt8.

source
Metal.simd_shuffle_xor — Function
simd_shuffle_xor(data::T, mask::Integer)

Return data of the thread in the SIMD-group whose 0-based index is the bitwise xor of the caller's 0-based index and mask, which has to be the same for all threads. T can be Float32, Float16, Int32, UInt32, Int16, UInt16, Int8 or UInt8.

source
Metal.simd_shuffle_down — Function
simd_shuffle_down(data::T, delta::Integer)

Return data of the thread delta lanes higher in the SIMD-group. The last delta threads, which have no such thread, get their own data. delta has to be the same for all threads. T can be Float32, Float16, Int32, UInt32, Int16, UInt16, Int8 or UInt8.

source
Metal.simd_shuffle_up — Function
simd_shuffle_up(data::T, delta::Integer)

Return data of the thread delta lanes lower in the SIMD-group. The first delta threads, which have no such thread, get their own data. delta has to be the same for all threads. T can be Float32, Float16, Int32, UInt32, Int16, UInt16, Int8 or UInt8.

source
Metal.simd_shuffle_and_fill_down — Function
simd_shuffle_and_fill_down(data::T, filling_data::T, delta::Integer,
                           [modulo::Integer=threads_per_simdgroup()])

Like simd_shuffle_down, but the last delta threads get filling_data of the first delta threads instead. With modulo, the SIMD-group is split into vectors of that many threads (2, 4, 8, 16 or 32), each shuffled on its own. delta and modulo have to be the same for all threads. T can be Float32, Float16, Int32, UInt32, Int16, UInt16, Int8 or UInt8.

source
Metal.simd_shuffle_and_fill_up — Function
simd_shuffle_and_fill_up(data::T, filling_data::T, delta::Integer,
                         [modulo::Integer=threads_per_simdgroup()])

Like simd_shuffle_up, but the first delta threads get filling_data of the last delta threads instead. With modulo, the SIMD-group is split into vectors of that many threads (2, 4, 8, 16 or 32), each shuffled on its own. delta and modulo have to be the same for all threads. T can be Float32, Float16, Int32, UInt32, Int16, UInt16, Int8 or UInt8.

source
Metal.simd_ballot — Function
simd_ballot(predicate::Bool)::UInt64

Return a bitmask in which bit i-1 is set if predicate is true for the active thread with index i in the SIMD-group. Bits of inactive threads are zero, and bits past the end of the SIMD-group are undefined: mask them off before, e.g., count_ones.

source
Metal.simd_vote_all — Function
simd_vote_all(bitmask::UInt64)::Bool

Return whether the bits of all active threads of the SIMD-group are set in bitmask, e.g., as returned by simd_ballot.

source
Metal.simd_vote_any — Function
simd_vote_any(bitmask::UInt64)::Bool

Return whether the bit of any active thread of the SIMD-group is set in bitmask, e.g., as returned by simd_ballot.

source

Quad-group functions

Functions that exchange data between the four threads of a quad-group.

Metal.quad_shuffle — Function
quad_shuffle(data::T, lane::Integer)

Return data of the thread in the quad-group whose 1-based index is lane (see thread_index_in_quadgroup). Different threads can read from different lanes. T can be Float32, Float16, Int32, UInt32, Int16, UInt16, Int8 or UInt8.

source
Metal.quad_shuffle_xor — Function
quad_shuffle_xor(data::T, mask::Integer)

Return data of the thread in the quad-group whose 0-based index is the bitwise xor of the caller's 0-based index and mask, which has to be the same for all threads. T can be Float32, Float16, Int32, UInt32, Int16, UInt16, Int8 or UInt8.

source
Metal.quad_shuffle_down — Function
quad_shuffle_down(data::T, delta::Integer)

Return data of the thread delta lanes higher in the quad-group. The last delta threads, which have no such thread, get their own data. delta has to be the same for all threads. T can be Float32, Float16, Int32, UInt32, Int16, UInt16, Int8 or UInt8.

source
Metal.quad_shuffle_up — Function
quad_shuffle_up(data::T, delta::Integer)

Return data of the thread delta lanes lower in the quad-group. The first delta threads, which have no such thread, get their own data. delta has to be the same for all threads. T can be Float32, Float16, Int32, UInt32, Int16, UInt16, Int8 or UInt8.

source
Metal.quad_shuffle_and_fill_down — Function
quad_shuffle_and_fill_down(data::T, filling_data::T, delta::Integer,
                           [modulo::Integer=4])

Like quad_shuffle_down, but the last delta threads get filling_data of the first delta threads instead. With modulo, the quad-group is split into vectors of that many threads (2 or 4), each shuffled on its own. delta and modulo have to be the same for all threads. T can be Float32, Float16, Int32, UInt32, Int16, UInt16, Int8 or UInt8.

source
Metal.quad_shuffle_and_fill_up — Function
quad_shuffle_and_fill_up(data::T, filling_data::T, delta::Integer,
                         [modulo::Integer=4])

Like quad_shuffle_up, but the first delta threads get filling_data of the last delta threads instead. With modulo, the quad-group is split into vectors of that many threads (2 or 4), each shuffled on its own. delta and modulo have to be the same for all threads. T can be Float32, Float16, Int32, UInt32, Int16, UInt16, Int8 or UInt8.

source
Metal.quad_ballot — Function
quad_ballot(predicate::Bool)::UInt64

Return a bitmask in which bit i-1 is set if predicate is true for the active thread with index i in the quad-group. Bits of inactive threads are zero, and bits past the end of the quad-group are undefined: mask them off before, e.g., count_ones.

source
Metal.quad_vote_all — Function
quad_vote_all(bitmask::UInt64)::Bool

Return whether the bits of all active threads of the quad-group are set in bitmask, e.g., as returned by quad_ballot.

source
Metal.quad_vote_any — Function
quad_vote_any(bitmask::UInt64)::Bool

Return whether the bit of any active thread of the quad-group is set in bitmask, e.g., as returned by quad_ballot.

source

SIMD-group matrices

8×8 matrices distributed over the threads of a SIMD-group, multiplied by its threads together.

Metal.MtlSimdgroupMatrix — Type
MtlSimdgroupMatrix{T,R,C}

Typed wrapper around a SIMD-group matrix fragment. T is the element type (Float16, Float32, or BFloat16); R and C are the matrix dimensions. Only the 8×8 shape is supported by current Apple GPUs.

The fragment data is distributed across the 32 lanes of a SIMD-group; the per-lane element layout is implementation-defined and elements cannot be accessed directly. To inspect or modify individual entries, store the matrix to device or threadgroup memory first.

Construct via simdgroup_load, zero or the explicit fill constructor MtlSimdgroupMatrix{T,8,8}(val::T).

source
Metal.simdgroup_load — Function
simdgroup_load(data::MtlDeviceArray{T}, matrix_origin=(1, 1), Val(transpose)=Val(true))

Loads data from device or threadgroup memory into an 8x8 SIMD-group matrix and returns it. T must be either Float16, Float32, or BFloat16.

Arguments

  • matrix_origin::NTuple{2, Int64}=(1, 1): origin in the source memory to load from.
  • Val(transpose)::Val{Bool}=Val(true): whether to transpose the loaded matrix. The default Val(true) treats the column-major source as a regular (non-transposed) matrix; pass Val(false) to load the transpose.
source
Metal.simdgroup_store — Function
simdgroup_store(src, dest::MtlDeviceArray{T}, matrix_origin=(1, 1), Val(transpose)=Val(true))

Stores data from an 8x8 SIMD-group matrix into device or threadgroup memory. T must be either Float16, Float32, or BFloat16.

Arguments

  • matrix_origin::NTuple{2, Int64}=(1, 1): origin in the destination memory to store to.
  • Val(transpose)::Val{Bool}=Val(true): whether to transpose the matrix on store. The default Val(true) matches the column-major convention; pass Val(false) to store the transpose.
source

Tensor operations

Matrix multiplication with the Metal 4 tensor operations, on GPUs of the Metal 4 family with macOS 26.

Metal.MtlInlineTensor — Type
MtlInlineTensor{T, R, AS}

A tensor view of an MtlDeviceArray or MtlThreadGroupArray, for the Metal 4 tensor operations (MSL's tensor_inline), which need a GPU in the Metal 4 family and macOS 26. T is the element type, R the rank and AS the address space of the array.

Extents and strides follow Julia's column-major convention, matching the array: the first dimension is contiguous. MtlInlineTensor(A) views all of A, and MtlInlineTensor(A, dims) and MtlInlineTensor(A, dims, strides) set an explicit shape. view(t, origin, dims) slices a tile, with a 1-based origin.

source
Metal.TensorOpsMatmul2D — Type
TensorOpsMatmul2D{DESC, NSIMD}

A matrix multiplication with Metal 4 tensor operations (MSL's tensor_ops::matmul2d), on MtlInlineTensors. DESC is a matmul2d_descriptor, and NSIMD the number of SIMD-groups that execute the operation together (execution_simdgroups<N> in MSL).

Construct it with TensorOpsMatmul2D(desc, Val(NSIMD)), and call it to compute C = A*B (or C += A*B with mode = matmul2d_multiply_accumulate):

op = TensorOpsMatmul2D(matmul2d_descriptor(64, 32), Val(4))op(A, B, C)            # C[64, 32] = A[64, K] * B[K, 32]
source
Metal.matmul2d_descriptor — Type
matmul2d_descriptor(M, N, [K]; transpose_left=false, transpose_right=false,
                    relaxed_precision=false, mode=matmul2d_multiply)

Describe a matrix multiplication C[M, N] = A[M, K] * B[K, N] for TensorOpsMatmul2D, like MSL's matmul2d_descriptor. K defaults to -1, meaning it is taken from the operands at run time. transpose_left and transpose_right transpose A and B; mode is a Matmul2DMode.

For an outer K-loop where each iteration accumulates a partial product into the destination, set mode = matmul2d_multiply_accumulate and zero the destination before the loop. A typical pattern:

# C[M, N] = A[M, K] * B[K, N], accumulated over K-tiles of width TileKop = TensorOpsMatmul2D{matmul2d_descriptor(M, N, TileK;                                           mode = matmul2d_multiply_accumulate),                       Int32(NSIMD)}()for s in 0:(nslices - 1)    k0 = Int32(s) * Int32(TileK)    sA = view(tA, (Int32(1), k0 + Int32(1)), (Int32(M), Int32(TileK)))    sB = view(tB, (k0 + Int32(1), Int32(1)), (Int32(TileK), Int32(N)))    op(sA, sB, tC)end

Limitation: the loop trip count needs to be kept dynamic (not a compile-time constant) to avoid crashing Apple's back-end.

source

Atomics

Atomic operations on device and threadgroup memory. The low-level functions follow MSL's atomic_*_explicit functions, taking pointers and optional memory orders; Metal.@atomic offers a higher-level interface on arrays.

Metal.@atomic — Macro
Metal.@atomic a[I] = op(a[I], val)
@atomic a[I] ...= val

Atomically perform a sequence of operations that loads an array element a[I], performs the operation op on that value and a second value val, and writes the result back to the array. This sequence can be written out as a regular assignment, in which case the same array element should be used in the left and right hand side of the assignment, or as an in-place application of a known operator. In both cases, the array reference should be pure and not induce any side-effects.

@atomic a[I] atomically loads an element. The element type and memory have to be supported by the low-level functions such as Metal.atomic_fetch_add_explicit: operations without a native atomic, e.g., *, are implemented with a compare-exchange loop. All operations are relaxed (Metal.memory_order_relaxed); use the low-level functions for other memory orders.

This interface is experimental, and might change without deprecation.

function histogram(counts, data)    i = thread_position_in_grid().x    if i <= length(data)        Metal.@atomic counts[data[i]] += 1    end    returnend
source
Metal.atomic_load_explicit — Function
atomic_load_explicit(ptr::LLVMPtr{T}, [order, [flags]]) -> T

Atomically load the value at ptr, which can be an Int32, UInt32 or Float32 in device or threadgroup memory.

The order is a memory_order (or a Val of one), Metal.memory_order_relaxed by default, and has to be a compile-time constant, as in MSL. As in MSL, the operation synchronizes with the other threads on the device for device memory, and with those in the threadgroup for threadgroup memory. An ordered operation orders accesses to both device and threadgroup memory; before Metal 4.1, it is implemented with fences. A load only uses the acquire part of an order, and a store only its release part.

The optional flags (MemoryFlags), also a compile-time constant, select the memory an ordered operation orders instead: e.g., only device memory (Metal.MemoryFlagDevice), or also texture memory (Metal.MemoryFlagDevice | Metal.MemoryFlagThreadGroup | Metal.MemoryFlagTexture).

As in MSL, a relaxed load of device memory is only guaranteed to eventually observe stores from threads in the same threadgroup: repeating it, e.g., in a loop waiting for a flag, may never observe a store from another threadgroup. To wait for another threadgroup, use an acquire load (memory_order_acquire), or a relaxed load at device scope through UnsafeAtomics (UnsafeAtomics.load(ptr, UnsafeAtomics.monotonic)).

source
Metal.atomic_store_explicit — Function
atomic_store_explicit(ptr::LLVMPtr{T}, val::T, [order, [flags]])

Atomically store val at ptr, which can be an Int32, UInt32 or Float32 in device or threadgroup memory.

The order is a memory_order (or a Val of one), Metal.memory_order_relaxed by default, and has to be a compile-time constant, as in MSL. As in MSL, the operation synchronizes with the other threads on the device for device memory, and with those in the threadgroup for threadgroup memory. An ordered operation orders accesses to both device and threadgroup memory; before Metal 4.1, it is implemented with fences. A load only uses the acquire part of an order, and a store only its release part.

The optional flags (MemoryFlags), also a compile-time constant, select the memory an ordered operation orders instead: e.g., only device memory (Metal.MemoryFlagDevice), or also texture memory (Metal.MemoryFlagDevice | Metal.MemoryFlagThreadGroup | Metal.MemoryFlagTexture).

source
Metal.atomic_exchange_explicit — Function
atomic_exchange_explicit(ptr::LLVMPtr{T}, val::T, [order, [flags]]) -> T

Atomically replace the value at ptr by val, and return the old value. T can be Int32, UInt32 or Float32, in device or threadgroup memory.

The order is a memory_order (or a Val of one), Metal.memory_order_relaxed by default, and has to be a compile-time constant, as in MSL. As in MSL, the operation synchronizes with the other threads on the device for device memory, and with those in the threadgroup for threadgroup memory. An ordered operation orders accesses to both device and threadgroup memory; before Metal 4.1, it is implemented with fences. A load only uses the acquire part of an order, and a store only its release part.

The optional flags (MemoryFlags), also a compile-time constant, select the memory an ordered operation orders instead: e.g., only device memory (Metal.MemoryFlagDevice), or also texture memory (Metal.MemoryFlagDevice | Metal.MemoryFlagThreadGroup | Metal.MemoryFlagTexture).

source
Metal.atomic_compare_exchange_weak_explicit — Function
atomic_compare_exchange_weak_explicit(ptr::LLVMPtr{T}, expected::T, desired::T,
                                      [success_order, failure_order, [flags]]) -> T

Atomically replace the value at ptr by desired if it equals expected, and return the value that was loaded. Unlike in MSL, this function doesn't return whether the exchange succeeded: it did if the returned value equals expected. As a weak compare-exchange, it can fail spuriously. T can be Int32, UInt32 or Float32, in device or threadgroup memory.

The success_order applies when the exchange succeeds, the failure_order (which cannot release) when it fails.

The order is a memory_order (or a Val of one), Metal.memory_order_relaxed by default, and has to be a compile-time constant, as in MSL. As in MSL, the operation synchronizes with the other threads on the device for device memory, and with those in the threadgroup for threadgroup memory. An ordered operation orders accesses to both device and threadgroup memory; before Metal 4.1, it is implemented with fences. A load only uses the acquire part of an order, and a store only its release part.

The optional flags (MemoryFlags), also a compile-time constant, select the memory an ordered operation orders instead: e.g., only device memory (Metal.MemoryFlagDevice), or also texture memory (Metal.MemoryFlagDevice | Metal.MemoryFlagThreadGroup | Metal.MemoryFlagTexture).

source
Metal.atomic_fetch_add_explicit — Function
atomic_fetch_add_explicit(ptr::LLVMPtr{T}, val::T, [order, [flags]]) -> T

Atomically add val to the value at ptr, and return the old value. T can be Int32, UInt32 or Float32, in device or threadgroup memory.

The order is a memory_order (or a Val of one), Metal.memory_order_relaxed by default, and has to be a compile-time constant, as in MSL. As in MSL, the operation synchronizes with the other threads on the device for device memory, and with those in the threadgroup for threadgroup memory. An ordered operation orders accesses to both device and threadgroup memory; before Metal 4.1, it is implemented with fences. A load only uses the acquire part of an order, and a store only its release part.

The optional flags (MemoryFlags), also a compile-time constant, select the memory an ordered operation orders instead: e.g., only device memory (Metal.MemoryFlagDevice), or also texture memory (Metal.MemoryFlagDevice | Metal.MemoryFlagThreadGroup | Metal.MemoryFlagTexture).

source
Metal.atomic_fetch_sub_explicit — Function
atomic_fetch_sub_explicit(ptr::LLVMPtr{T}, val::T, [order, [flags]]) -> T

Atomically subtract val from the value at ptr, and return the old value. T can be Int32, UInt32 or Float32, in device or threadgroup memory.

The order is a memory_order (or a Val of one), Metal.memory_order_relaxed by default, and has to be a compile-time constant, as in MSL. As in MSL, the operation synchronizes with the other threads on the device for device memory, and with those in the threadgroup for threadgroup memory. An ordered operation orders accesses to both device and threadgroup memory; before Metal 4.1, it is implemented with fences. A load only uses the acquire part of an order, and a store only its release part.

The optional flags (MemoryFlags), also a compile-time constant, select the memory an ordered operation orders instead: e.g., only device memory (Metal.MemoryFlagDevice), or also texture memory (Metal.MemoryFlagDevice | Metal.MemoryFlagThreadGroup | Metal.MemoryFlagTexture).

source
Metal.atomic_fetch_min_explicit — Function
atomic_fetch_min_explicit(ptr::LLVMPtr{T}, val::T, [order, [flags]]) -> T

Atomically replace the value at ptr by its minimum with val, and return the old value. T can be Int32 or UInt32, in device or threadgroup memory.

The order is a memory_order (or a Val of one), Metal.memory_order_relaxed by default, and has to be a compile-time constant, as in MSL. As in MSL, the operation synchronizes with the other threads on the device for device memory, and with those in the threadgroup for threadgroup memory. An ordered operation orders accesses to both device and threadgroup memory; before Metal 4.1, it is implemented with fences. A load only uses the acquire part of an order, and a store only its release part.

The optional flags (MemoryFlags), also a compile-time constant, select the memory an ordered operation orders instead: e.g., only device memory (Metal.MemoryFlagDevice), or also texture memory (Metal.MemoryFlagDevice | Metal.MemoryFlagThreadGroup | Metal.MemoryFlagTexture).

source
Metal.atomic_fetch_max_explicit — Function
atomic_fetch_max_explicit(ptr::LLVMPtr{T}, val::T, [order, [flags]]) -> T

Atomically replace the value at ptr by its maximum with val, and return the old value. T can be Int32 or UInt32, in device or threadgroup memory.

The order is a memory_order (or a Val of one), Metal.memory_order_relaxed by default, and has to be a compile-time constant, as in MSL. As in MSL, the operation synchronizes with the other threads on the device for device memory, and with those in the threadgroup for threadgroup memory. An ordered operation orders accesses to both device and threadgroup memory; before Metal 4.1, it is implemented with fences. A load only uses the acquire part of an order, and a store only its release part.

The optional flags (MemoryFlags), also a compile-time constant, select the memory an ordered operation orders instead: e.g., only device memory (Metal.MemoryFlagDevice), or also texture memory (Metal.MemoryFlagDevice | Metal.MemoryFlagThreadGroup | Metal.MemoryFlagTexture).

source
Metal.atomic_fetch_and_explicit — Function
atomic_fetch_and_explicit(ptr::LLVMPtr{T}, val::T, [order, [flags]]) -> T

Atomically replace the value at ptr by its bitwise and with val, and return the old value. T can be Int32 or UInt32, in device or threadgroup memory.

The order is a memory_order (or a Val of one), Metal.memory_order_relaxed by default, and has to be a compile-time constant, as in MSL. As in MSL, the operation synchronizes with the other threads on the device for device memory, and with those in the threadgroup for threadgroup memory. An ordered operation orders accesses to both device and threadgroup memory; before Metal 4.1, it is implemented with fences. A load only uses the acquire part of an order, and a store only its release part.

The optional flags (MemoryFlags), also a compile-time constant, select the memory an ordered operation orders instead: e.g., only device memory (Metal.MemoryFlagDevice), or also texture memory (Metal.MemoryFlagDevice | Metal.MemoryFlagThreadGroup | Metal.MemoryFlagTexture).

source
Metal.atomic_fetch_or_explicit — Function
atomic_fetch_or_explicit(ptr::LLVMPtr{T}, val::T, [order, [flags]]) -> T

Atomically replace the value at ptr by its bitwise or with val, and return the old value. T can be Int32 or UInt32, in device or threadgroup memory.

The order is a memory_order (or a Val of one), Metal.memory_order_relaxed by default, and has to be a compile-time constant, as in MSL. As in MSL, the operation synchronizes with the other threads on the device for device memory, and with those in the threadgroup for threadgroup memory. An ordered operation orders accesses to both device and threadgroup memory; before Metal 4.1, it is implemented with fences. A load only uses the acquire part of an order, and a store only its release part.

The optional flags (MemoryFlags), also a compile-time constant, select the memory an ordered operation orders instead: e.g., only device memory (Metal.MemoryFlagDevice), or also texture memory (Metal.MemoryFlagDevice | Metal.MemoryFlagThreadGroup | Metal.MemoryFlagTexture).

source
Metal.atomic_fetch_xor_explicit — Function
atomic_fetch_xor_explicit(ptr::LLVMPtr{T}, val::T, [order, [flags]]) -> T

Atomically replace the value at ptr by its bitwise xor with val, and return the old value. T can be Int32 or UInt32, in device or threadgroup memory.

The order is a memory_order (or a Val of one), Metal.memory_order_relaxed by default, and has to be a compile-time constant, as in MSL. As in MSL, the operation synchronizes with the other threads on the device for device memory, and with those in the threadgroup for threadgroup memory. An ordered operation orders accesses to both device and threadgroup memory; before Metal 4.1, it is implemented with fences. A load only uses the acquire part of an order, and a store only its release part.

The optional flags (MemoryFlags), also a compile-time constant, select the memory an ordered operation orders instead: e.g., only device memory (Metal.MemoryFlagDevice), or also texture memory (Metal.MemoryFlagDevice | Metal.MemoryFlagThreadGroup | Metal.MemoryFlagTexture).

source
Metal.atomic_min_explicit — Function
atomic_min_explicit(ptr::LLVMPtr{UInt64,AS.Device}, val::UInt64, [order, [flags]])

Atomically replace the UInt64 in device memory at ptr by its minimum with val. Unlike the 32-bit operations, this doesn't return the old value, as Metal only supports the non-fetching form. It needs an Apple8 GPU or newer.

source
Metal.atomic_max_explicit — Function
atomic_max_explicit(ptr::LLVMPtr{UInt64,AS.Device}, val::UInt64, [order, [flags]])

Atomically replace the UInt64 in device memory at ptr by its maximum with val. Unlike the 32-bit operations, this doesn't return the old value, as Metal only supports the non-fetching form. It needs an Apple8 GPU or newer.

source
Metal.atomic_fetch_op_explicit — Function
atomic_fetch_op_explicit(ptr::LLVMPtr{T}, op, val, [order, [flags]]) -> T

Atomically replace the value at ptr by op(old, val), and return the old value old. Implemented with a compare-exchange loop, this supports any operation, but only the types and address spaces Metal.atomic_compare_exchange_weak_explicit does.

The order is a memory_order (or a Val of one), Metal.memory_order_relaxed by default, and has to be a compile-time constant, as in MSL. As in MSL, the operation synchronizes with the other threads on the device for device memory, and with those in the threadgroup for threadgroup memory. An ordered operation orders accesses to both device and threadgroup memory; before Metal 4.1, it is implemented with fences. A load only uses the acquire part of an order, and a store only its release part.

The optional flags (MemoryFlags), also a compile-time constant, select the memory an ordered operation orders instead: e.g., only device memory (Metal.MemoryFlagDevice), or also texture memory (Metal.MemoryFlagDevice | Metal.MemoryFlagThreadGroup | Metal.MemoryFlagTexture).

source

Printing

Metal.@mtlprintf — Macro
@mtlprintf(fmt::String, args...)

Print a formatted string from a kernel, with C printf formatting. The format string has to be a literal. Arguments are passed as C variadic arguments, so use the length modifier of their type (e.g., %d for Int32, %ld for Int64, %u for UInt32, %f for Float32); the %n and %s conversions are not supported.

Metal collects the output in a log buffer, which Metal.jl prints to the host's standard output when the kernel's command buffer completes. When the log buffer is full, further messages are dropped. Kernels that print are submitted in a command buffer of their own, bypassing command batching, and cannot run on virtualized GPUs or during a GPU frame capture.

@mtlprintf("thread %u: %f", thread_position_in_grid().x, 1f0)
source
Metal.@mtlprint — Macro
@mtlprint(xs...)
@mtlprintln(xs...)

Print a textual representation of values xs from a kernel, followed by a newline for @mtlprintln. This builds on @mtlprintf, choosing the format specifiers for you, and supports 16-, 32- and 64-bit integers, Float32, Bool, Cchar, pointers, and tuples of those. For other output, use @mtlprintf directly.

Limited string interpolation is also possible:

    @mtlprint("Hello, World ", 42, "\n")    @mtlprint "Hello, World $(42)\n"
source
Metal.@mtlprintln — Macro
@mtlprint(xs...)
@mtlprintln(xs...)

Print a textual representation of values xs from a kernel, followed by a newline for @mtlprintln. This builds on @mtlprintf, choosing the format specifiers for you, and supports 16-, 32- and 64-bit integers, Float32, Bool, Cchar, pointers, and tuples of those. For other output, use @mtlprintf directly.

Limited string interpolation is also possible:

    @mtlprint("Hello, World ", 42, "\n")    @mtlprint "Hello, World $(42)\n"
source

Random numbers

In kernels, rand(), rand(T), randn() and randn(T) from the Random standard library draw from a Philox2x32 generator, which every launch seeds with a new random seed. Call Random.seed!(seed) in the kernel, from every SIMD-group, for reproducible numbers.