diff --git a/dev/.documenter-siteinfo.json b/dev/.documenter-siteinfo.json index 85e420b4..fc4896f6 100644 --- a/dev/.documenter-siteinfo.json +++ b/dev/.documenter-siteinfo.json @@ -1 +1 @@ -{"documenter":{"julia_version":"1.10.4","generation_timestamp":"2024-08-23T07:50:36","documenter_version":"1.5.0"}} \ No newline at end of file +{"documenter":{"julia_version":"1.10.5","generation_timestamp":"2024-08-29T13:41:43","documenter_version":"1.5.0"}} \ No newline at end of file diff --git a/dev/api/array/index.html b/dev/api/array/index.html index 3909a455..7687b073 100644 --- a/dev/api/array/index.html +++ b/dev/api/array/index.html @@ -14,4 +14,4 @@ 3-element MtlVector{Int64, Metal.PrivateStorage}: 1 2 - 3source
Metal.MtlArray — Type
MtlArray{T,N,S} <: AbstractGPUArray{T,N}

N-dimensional Metal array with storage mode S and elements of type T.

S can be Metal.PrivateStorage (default), Metal.SharedStorage, or Metal.ManagedStorage.

See the Array Programming section of the Metal.jl docs for more details.

source
Metal.MtlVector — Type
MtlVector{T,S} <: AbstractGPUVector{T}

One-dimensional array with elements of type T for use with Apple Metal-compatible GPUs. Alias for MtlArray{T,1,S}.

See also Vector(@ref), and the Array Programming section of the Metal.jl docs for more details.

source
Metal.MtlMatrix — Type
MtlMatrix{T,S} <: AbstractGPUMatrix{T}

Two-dimensional array with elements of type T for use with Apple Metal-compatible GPUs. Alias for MtlArray{T,2,S}.

See also Matrix(@ref), and the Array Programming section of the Metal.jl docs for more details.

source
Metal.MtlVecOrMat — Type
MtlVecOrMat{T,S}

Union type of MtlVector{T,S} and MtlMatrix{T,S} which allows functions to accept either an MtlMatrix or an MtlVector.

See also VecOrMat(@ref) for examples.

source

Storage modes

The Metal API has various storage modes that dictate how a resource can be accessed. MtlArrays are Metal.PrivateStorage by default, but they can also be Metal.SharedStorage or Metal.ManagedStorage. For more information on storage modes, see the official Metal documentation.

Metal.MTL.PrivateStorage — Type
struct Metal.PrivateStorage <: MTL.StorageMode

Used to indicate that the resource is stored using MTLStorageModePrivate in memory.

For more information on Metal storage modes, refer to the official Metal documentation.

See also Metal.SharedStorage and Metal.ManagedStorage.

source
Metal.MTL.SharedStorage — Type
struct Metal.SharedStorage <: MTL.StorageMode

Used to indicate that the resource is stored using MTLStorageModeShared in memory.

For more information on Metal storage modes, refer to the official Metal documentation.

See also Metal.PrivateStorage and Metal.ManagedStorage.

source
Metal.MTL.ManagedStorage — Type
struct Metal.ManagedStorage <: MTL.StorageMode

Used to indicate that the resource is stored using MTLStorageModeManaged in memory.

For more information on Metal storage modes, refer to the official Metal documentation.

See also Metal.SharedStorage and Metal.PrivateStorage.

source

There also exist the following convenience functions to check if an MtlArray is using a specific storage mode:

Metal.is_private — Function
is_private(A::MtlArray) -> Bool

Returns true if A has storage mode Metal.PrivateStorage.

See also is_shared and is_managed.

source
Metal.is_shared — Function
is_shared(A::MtlArray) -> Bool

Returns true if A has storage mode Metal.SharedStorage.

See also is_private and is_managed.

source
Metal.is_managed — Function
is_managed(A::MtlArray) -> Bool

Returns true if A has storage mode Metal.ManagedStorage.

See also is_shared and is_private.

source
+ 3source
Metal.MtlArray — Type
MtlArray{T,N,S} <: AbstractGPUArray{T,N}

N-dimensional Metal array with storage mode S and elements of type T.

S can be Metal.PrivateStorage (default), Metal.SharedStorage, or Metal.ManagedStorage.

See the Array Programming section of the Metal.jl docs for more details.

source
Metal.MtlVector — Type
MtlVector{T,S} <: AbstractGPUVector{T}

One-dimensional array with elements of type T for use with Apple Metal-compatible GPUs. Alias for MtlArray{T,1,S}.

See also Vector(@ref), and the Array Programming section of the Metal.jl docs for more details.

source
Metal.MtlMatrix — Type
MtlMatrix{T,S} <: AbstractGPUMatrix{T}

Two-dimensional array with elements of type T for use with Apple Metal-compatible GPUs. Alias for MtlArray{T,2,S}.

See also Matrix(@ref), and the Array Programming section of the Metal.jl docs for more details.

source
Metal.MtlVecOrMat — Type
MtlVecOrMat{T,S}

Union type of MtlVector{T,S} and MtlMatrix{T,S} which allows functions to accept either an MtlMatrix or an MtlVector.

See also VecOrMat(@ref) for examples.

source

Storage modes

The Metal API has various storage modes that dictate how a resource can be accessed. MtlArrays are Metal.PrivateStorage by default, but they can also be Metal.SharedStorage or Metal.ManagedStorage. For more information on storage modes, see the official Metal documentation.

Metal.MTL.PrivateStorage — Type
struct Metal.PrivateStorage <: MTL.StorageMode

Used to indicate that the resource is stored using MTLStorageModePrivate in memory.

For more information on Metal storage modes, refer to the official Metal documentation.

See also Metal.SharedStorage and Metal.ManagedStorage.

source
Metal.MTL.SharedStorage — Type
struct Metal.SharedStorage <: MTL.StorageMode

Used to indicate that the resource is stored using MTLStorageModeShared in memory.

For more information on Metal storage modes, refer to the official Metal documentation.

See also Metal.PrivateStorage and Metal.ManagedStorage.

source
Metal.MTL.ManagedStorage — Type
struct Metal.ManagedStorage <: MTL.StorageMode

Used to indicate that the resource is stored using MTLStorageModeManaged in memory.

For more information on Metal storage modes, refer to the official Metal documentation.

See also Metal.SharedStorage and Metal.PrivateStorage.

source

There also exist the following convenience functions to check if an MtlArray is using a specific storage mode:

Metal.is_private — Function
is_private(A::MtlArray) -> Bool

Returns true if A has storage mode Metal.PrivateStorage.

See also is_shared and is_managed.

source
Metal.is_shared — Function
is_shared(A::MtlArray) -> Bool

Returns true if A has storage mode Metal.SharedStorage.

See also is_private and is_managed.

source
Metal.is_managed — Function
is_managed(A::MtlArray) -> Bool

Returns true if A has storage mode Metal.ManagedStorage.

See also is_shared and is_private.

source
diff --git a/dev/api/compiler/index.html b/dev/api/compiler/index.html index ba67a885..ec7c5244 100644 --- a/dev/api/compiler/index.html +++ b/dev/api/compiler/index.html @@ -1,8 +1,8 @@ -Compiler · Metal.jl

Compiler

Execution

The main entry-point to the compiler is the @metal macro:

Metal.@metal — Macro
@metal threads=... groups=... [kwargs...] func(args...)

High-level interface for executing code on a GPU.

The @metal macro should prefix a call, with func a callable function or object that should return nothing. It will be compiled to a Metal function upon first use, and to a certain extent arguments will be converted and managed automatically using mtlconvert. Finally, a call to mtlcall is performed, creating a command buffer in the current global command queue then committing it.

There is one supported keyword argument that influences the behavior of @metal:

  • launch: whether to launch this kernel, defaults to true. If false the returned kernel object should be launched by calling it and passing arguments again.
  • name: the name of the kernel in the generated code. Defaults to an automatically- generated name.
  • queue: the command queue to use for this kernel. Defaults to the global command queue.
source

If needed, you can use a lower-level API that lets you inspect the compiler kernel:

Metal.mtlconvert — Function

mtlconvert(x, [cce])

This function is called for every argument to be passed to a kernel, allowing it to be converted to a GPU-friendly format. By default, the function does nothing and returns the input object x as-is.

Do not add methods to this function, but instead extend the underlying Adapt.jl package and register methods for the the Metal.Adaptor type.

source
Metal.mtlfunction — Function
mtlfunction(f, tt=Tuple{}; kwargs...)

Low-level interface to compile a function invocation for the currently-active GPU, returning a callable kernel object. For a higher-level interface, use @metal.

The output of this function is automatically cached, i.e. you can simply call mtlfunction in a hot path without degrading performance. New code will be generated automatically when the function changes, or when different types or keyword arguments are provided.

source

Reflection

If you want to inspect generated code, you can use macros that resemble functionality from the InteractiveUtils standard library:

@device_code_lowered
+Compiler · Metal.jl
Compiler
Execution
The main entry-point to the compiler is the @metal macro:
Metal.@metal — Macro
@metal threads=... groups=... [kwargs...] func(args...)
High-level interface for executing code on a GPU.
The @metal macro should prefix a call, with func a callable function or object that should return nothing. It will be compiled to a Metal function upon first use, and to a certain extent arguments will be converted and managed automatically using mtlconvert. Finally, a call to mtlcall is performed, creating a command buffer in the current global command queue then committing it.
There is one supported keyword argument that influences the behavior of @metal:
  • launch: whether to launch this kernel, defaults to true. If false the returned kernel object should be launched by calling it and passing arguments again.
  • name: the name of the kernel in the generated code. Defaults to an automatically- generated name.
  • queue: the command queue to use for this kernel. Defaults to the global command queue.
source
If needed, you can use a lower-level API that lets you inspect the compiler kernel:
Metal.mtlconvert — Function
mtlconvert(x, [cce])
This function is called for every argument to be passed to a kernel, allowing it to be converted to a GPU-friendly format. By default, the function does nothing and returns the input object x as-is.
Do not add methods to this function, but instead extend the underlying Adapt.jl package and register methods for the the Metal.Adaptor type.
source
Metal.mtlfunction — Function
mtlfunction(f, tt=Tuple{}; kwargs...)
Low-level interface to compile a function invocation for the currently-active GPU, returning a callable kernel object. For a higher-level interface, use @metal.
The output of this function is automatically cached, i.e. you can simply call mtlfunction in a hot path without degrading performance. New code will be generated automatically when the function changes, or when different types or keyword arguments are provided.
source
Reflection
If you want to inspect generated code, you can use macros that resemble functionality from the InteractiveUtils standard library:
@device_code_lowered
 @device_code_typed
 @device_code_warntype
 @device_code_llvm
 @device_code_native
 @device_code_agx
-@device_code
For more information, please consult the GPUCompiler.jl documentation. code_agx is actually code_native:

+@device_code

For more information, please consult the GPUCompiler.jl documentation. code_agx is actually code_native:

diff --git a/dev/api/essentials/index.html b/dev/api/essentials/index.html index cbf1a8cd..31424d5f 100644 --- a/dev/api/essentials/index.html +++ b/dev/api/essentials/index.html @@ -1,2 +1,2 @@ -Essentials · Metal.jl

Essentials

Versions and Support

Global State

Metal.device! — Function
device!(dev::MTLDevice)

Sets the Metal GPU device associated with the current Julia task.

source
Metal.device — Function
device()::MTLDevice

Return the Metal GPU device associated with the current Julia task.

Since all M-series systems currently only externally show a single GPU, this function effectively returns the only system GPU.

source
device(<:MtlArray)

Get the Metal device for an MtlArray.

source
Metal.global_queue — Function
global_queue(dev::MTLDevice)::MTLCommandQueue

Return the Metal command queue associated with the current Julia thread.

source
Metal.synchronize — Function
synchronize(queue)

Wait for currently committed GPU work on this queue to finish.

Create a new MTLCommandBuffer from the global command queue, commit it to the queue, and simply wait for it to be completed. Since command buffers should execute in a First-In-First-Out manner, this synchronizes the GPU.

source
+Essentials · Metal.jl

Essentials

Versions and Support

Global State

Metal.device! — Function
device!(dev::MTLDevice)

Sets the Metal GPU device associated with the current Julia task.

source
Metal.device — Function
device()::MTLDevice

Return the Metal GPU device associated with the current Julia task.

Since all M-series systems currently only externally show a single GPU, this function effectively returns the only system GPU.

source
device(<:MtlArray)

Get the Metal device for an MtlArray.

source
Metal.global_queue — Function
global_queue(dev::MTLDevice)::MTLCommandQueue

Return the Metal command queue associated with the current Julia thread.

source
Metal.synchronize — Function
synchronize(queue)

Wait for currently committed GPU work on this queue to finish.

Create a new MTLCommandBuffer from the global command queue, commit it to the queue, and simply wait for it to be completed. Since command buffers should execute in a First-In-First-Out manner, this synchronizes the GPU.

source
diff --git a/dev/api/kernel/index.html b/dev/api/kernel/index.html index ed522e67..b16d5527 100644 --- a/dev/api/kernel/index.html +++ b/dev/api/kernel/index.html @@ -1,24 +1,24 @@ -Kernel programming · Metal.jl

Kernel programming

This section lists the package's public functionality that corresponds to special Metal functions for use in device code. For more information about these functions, please consult the Metal Shading Language specification.

This is made possible by interfacing with the Metal libraries by wrapping a subset of the ObjectiveC APIs using ObjectiveC.jl. These low-level wrappers are available in the MTL submodule exported by Metal.jl.

Indexing and dimensions

Metal.thread_position_in_grid_1d — Function
thread_position_in_grid_1d()::UInt32
+Kernel programming · Metal.jl
Kernel programming
This section lists the package's public functionality that corresponds to special Metal functions for use in device code. For more information about these functions, please consult the Metal Shading Language specification.
This is made possible by interfacing with the Metal libraries by wrapping a subset of the ObjectiveC APIs using ObjectiveC.jl. These low-level wrappers are available in the MTL submodule exported by Metal.jl.
Indexing and dimensions
Metal.thread_position_in_grid_1d — Function
thread_position_in_grid_1d()::UInt32
 thread_position_in_grid_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}
-thread_position_in_grid_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}
Return the current thread's position in an N-dimensional grid of threads.
source
Metal.thread_position_in_threadgroup_1d — Function
thread_position_in_threadgroup_1d()::UInt32
+thread_position_in_grid_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}
Return the current thread's position in an N-dimensional grid of threads.
source
Metal.thread_position_in_threadgroup_1d — Function
thread_position_in_threadgroup_1d()::UInt32
 thread_position_in_threadgroup_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}
-thread_position_in_threadgroup_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}
Return the current thread's unique position within a threadgroup.
source
Metal.threadgroup_position_in_grid_1d — Function
threadgroup_position_in_grid_1d()::UInt32
+thread_position_in_threadgroup_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}
Return the current thread's unique position within a threadgroup.
source
Metal.threadgroup_position_in_grid_1d — Function
threadgroup_position_in_grid_1d()::UInt32
 threadgroup_position_in_grid_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}
-threadgroup_position_in_grid_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}
Return the current threadgroup's unique position within the grid.
source
Metal.threadgroups_per_grid_1d — Function
threadgroups_per_grid_1d()::UInt32
+threadgroup_position_in_grid_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}
Return the current threadgroup's unique position within the grid.
source
Metal.threadgroups_per_grid_1d — Function
threadgroups_per_grid_1d()::UInt32
 threadgroups_per_grid_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}
-threadgroups_per_grid_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}
Return the number of threadgroups per grid.
source
Metal.threads_per_grid_1d — Function
threads_per_grid_1d()::UInt32
+threadgroups_per_grid_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}
Return the number of threadgroups per grid.
source
Metal.threads_per_grid_1d — Function
threads_per_grid_1d()::UInt32
 threads_per_grid_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}
-threads_per_grid_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}
Return the grid size.
source
Metal.threads_per_threadgroup_1d — Function
threads_per_threadgroup_1d()::UInt32
 threads_per_threadgroup_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}
-threads_per_threadgroup_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}
Return the thread execution width of a threadgroup.
source
Metal.grid_size_1d — Function
grid_size_1d()::UInt32
+threads_per_threadgroup_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}
Return the thread execution width of a threadgroup.
source
Metal.grid_size_1d — Function
grid_size_1d()::UInt32
 grid_size_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}
-grid_size_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}
Return maximum size of the grid for threads that read per-thread stage-in data.
source
Metal.grid_origin_1d — Function
grid_origin_1d()::UInt32
+grid_size_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}
Return maximum size of the grid for threads that read per-thread stage-in data.
source
Metal.grid_origin_1d — Function
grid_origin_1d()::UInt32
 grid_origin_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}
-grid_origin_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}
Return the origin offset of the grid for threads that read per-thread stage-in data.
source
Device arrays
Metal.jl provides a primitive, lightweight array type to manage GPU data organized in an plain, dense fashion. This is the device-counterpart to the MtlArray, and implements (part of) the array interface as well as other functionality for use on the GPU:
Metal.MtlDeviceArray — Type
MtlDeviceArray(dims, ptr)
+grid_origin_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}
Return the origin offset of the grid for threads that read per-thread stage-in data.
source
Device arrays
Metal.jl provides a primitive, lightweight array type to manage GPU data organized in an plain, dense fashion. This is the device-counterpart to the MtlArray, and implements (part of) the array interface as well as other functionality for use on the GPU:
Metal.MtlDeviceArray — Type
MtlDeviceArray(dims, ptr)
 MtlDeviceArray{T}(dims, ptr)
 MtlDeviceArray{T,A}(dims, ptr)
-MtlDeviceArray{T,A,N}(dims, ptr)
Construct an N-dimensional dense Metal device array with element type T wrapping a pointer, where N is determined from the length of dims and T is determined from the type of ptr.
dims may be a single scalar, or a tuple of integers corresponding to the lengths in each dimension). If the rank N is supplied explicitly as in Array{T,N}(dims), then it must match the length of dims. The same applies to the element type T, which should match the type of the pointer ptr.
source
Metal.Const — Type
Const(A::MtlDeviceArray)
Mark a MtlDeviceArray as constant/read-only and to use the constant address space.
Warning
Experimental API. Subject to change without deprecation.
source
Shared memory
Metal.MtlThreadGroupArray — Function
MtlThreadGroupArray(::Type{T}, dims)
Create an array local to each threadgroup launched during kernel execution.
source
Synchronization
Metal.MemoryFlags — Type
MemoryFlags
Flags to set the memory synchronization behavior of threadgroup_barrier and simdgroup_barrier.
Possible values:
None: Set barriers to only act as an execution barrier and not apply a memory fence.
+MtlDeviceArray{T,A,N}(dims, ptr)
Construct an N-dimensional dense Metal device array with element type T wrapping a pointer, where N is determined from the length of dims and T is determined from the type of ptr.
dims may be a single scalar, or a tuple of integers corresponding to the lengths in each dimension). If the rank N is supplied explicitly as in Array{T,N}(dims), then it must match the length of dims. The same applies to the element type T, which should match the type of the pointer ptr.
source
Metal.Const — Type
Const(A::MtlDeviceArray)
Mark a MtlDeviceArray as constant/read-only and to use the constant address space.
Warning
Experimental API. Subject to change without deprecation.
source
Shared memory
Metal.MtlThreadGroupArray — Function
MtlThreadGroupArray(::Type{T}, dims)
Create an array local to each threadgroup launched during kernel execution.
source
Synchronization
Metal.MemoryFlags — Type
MemoryFlags
Flags to set the memory synchronization behavior of threadgroup_barrier and simdgroup_barrier.
Possible values:
None: Set barriers to only act as an execution barrier and not apply a memory fence.
 
 Device: Ensure the GPU correctly orders the memory operations to device memory
         for threads in the threadgroup or simdgroup.
@@ -30,4 +30,4 @@
         threads in a threadgroup or simdgroup for a texture with the read_write access qualifier.
 
 ThreadGroup_ImgBlock: Ensure the GPU correctly orders the memory operations to threadgroup imageblock memory
-        for threads in a threadgroup or simdgroup.
source
Metal.threadgroup_barrier — Function
threadgroup_barrier(flag::MemoryFlags=MemoryFlagNone)
Synchronize all threads in a threadgroup.
Possible flags that affect the memory synchronization behavior are found in MemoryFlags
source
Metal.simdgroup_barrier — Function
simdgroup_barrier(flag::MemoryFlags=MemoryFlagNone)
Synchronize all threads in a SIMD-group.
Possible flags that affect the memory synchronization behavior are found in MemoryFlags
source

+        for threads in a threadgroup or simdgroup.
source
Metal.threadgroup_barrier — Function
threadgroup_barrier(flag::MemoryFlags=MemoryFlagNone)

Synchronize all threads in a threadgroup.

Possible flags that affect the memory synchronization behavior are found in MemoryFlags

source
Metal.simdgroup_barrier — Function
simdgroup_barrier(flag::MemoryFlags=MemoryFlagNone)

Synchronize all threads in a SIMD-group.

Possible flags that affect the memory synchronization behavior are found in MemoryFlags

source
diff --git a/dev/api/mps/index.html b/dev/api/mps/index.html index 754bcd11..be6a7175 100644 --- a/dev/api/mps/index.html +++ b/dev/api/mps/index.html @@ -1,4 +1,4 @@ -Metal Performance Shaders · Metal.jl

Metal Performance Shaders

This section lists the package's public functionality that corresponds to the Metal Performance Shaders functions. For more information about these functions, or to see which functions have yet to be implemented in this package, please consult the Metal Performance Shaders Documentation.

Matrices and Vectors

Metal.MPS.MPSMatrix — Type
MPSMatrix(mat::MtlMatrix)

Metal matrix representation used in Performance Shaders.

Note that this results in a transposed view of the input, as Metal stores matrices row-major instead of column-major.

source
MPSMatrix(vec::MtlVector)

Metal matrix representation used in Performance Shaders.

Note that this results in a transposed view of the input, as Metal stores matrices row-major instead of column-major.

source
MPSMatrix(arr::MtlArray{T,3})

Metal batched matrix representation used in Performance Shaders.

Note that this results in a transposed view of the input, as Metal stores matrices row-major instead of column-major.

source

Matrix Arithmetic Operators

Metal.MPS.matmul! — Function
matMulMPS(a::MtlMatrix, b::MtlMatrix, c::MtlMatrix, alpha=1, beta=1,
-          transpose_left=false, transpose_right=false)

A MPSMatrixMultiplication kernel thay computes: c = alpha * op(a) * beta * op(b) + beta * C

This function should not typically be used. Rather, use the normal LinearAlgebra interface with any MtlArray and it should be accelerated using Metal Performance Shaders.

source
Metal.MPS.matvecmul! — Function
matvecmul!(c::MtlVector, a::MtlMatrix, b::MtlVector, alpha=1, beta=1, transpose=false)

A MPSMatrixVectorMultiplication kernel thay computes: c = alpha * op(a) * b + beta * c

This function should not typically be used. Rather, use the normal LinearAlgebra interface with any MtlArray and it should be accelerated using Metal Performance Shaders.

source
Metal.MPS.topk — Function
MPS.topk(A::MtlMatrix{T}, k) where {T<:MtlFloat}

Compute the top k values and their corresponding indices column-wise in a matrix A. Return the indices in I and the values in V.

k cannot be greater than 16.

Uses MPSMatrixFindTopK.

See also: topk!.

Warn

This interface is experimental, and might change without warning.

source
Metal.MPS.topk! — Function
MPS.topk!(A::MtlMatrix{T}, I::MtlMatrix{Int32}, V::MtlMatrix{T}, k)
-                                                 where {T<:MtlFloat}

Compute the top k values and their corresponding indices column-wise in a matrix A. Return the indices in I and the values in V.

k cannot be greater than 16.

Uses MPSMatrixFindTopK.

See also: topk.

Warn

This interface is experimental, and might change without warning.

source
+Metal Performance Shaders · Metal.jl

Metal Performance Shaders

This section lists the package's public functionality that corresponds to the Metal Performance Shaders functions. For more information about these functions, or to see which functions have yet to be implemented in this package, please consult the Metal Performance Shaders Documentation.

Matrices and Vectors

Metal.MPS.MPSMatrix — Type
MPSMatrix(mat::MtlMatrix)

Metal matrix representation used in Performance Shaders.

Note that this results in a transposed view of the input, as Metal stores matrices row-major instead of column-major.

source
MPSMatrix(vec::MtlVector)

Metal matrix representation used in Performance Shaders.

Note that this results in a transposed view of the input, as Metal stores matrices row-major instead of column-major.

source
MPSMatrix(arr::MtlArray{T,3})

Metal batched matrix representation used in Performance Shaders.

Note that this results in a transposed view of the input, as Metal stores matrices row-major instead of column-major.

source

Matrix Arithmetic Operators

Metal.MPS.matmul! — Function
matMulMPS(a::MtlMatrix, b::MtlMatrix, c::MtlMatrix, alpha=1, beta=1,
+          transpose_left=false, transpose_right=false)

A MPSMatrixMultiplication kernel thay computes: c = alpha * op(a) * beta * op(b) + beta * C

This function should not typically be used. Rather, use the normal LinearAlgebra interface with any MtlArray and it should be accelerated using Metal Performance Shaders.

source
Metal.MPS.matvecmul! — Function
matvecmul!(c::MtlVector, a::MtlMatrix, b::MtlVector, alpha=1, beta=1, transpose=false)

A MPSMatrixVectorMultiplication kernel thay computes: c = alpha * op(a) * b + beta * c

This function should not typically be used. Rather, use the normal LinearAlgebra interface with any MtlArray and it should be accelerated using Metal Performance Shaders.

source
Metal.MPS.topk — Function
MPS.topk(A::MtlMatrix{T}, k) where {T<:MtlFloat}

Compute the top k values and their corresponding indices column-wise in a matrix A. Return the indices in I and the values in V.

k cannot be greater than 16.

Uses MPSMatrixFindTopK.

See also: topk!.

Warn

This interface is experimental, and might change without warning.

source
Metal.MPS.topk! — Function
MPS.topk!(A::MtlMatrix{T}, I::MtlMatrix{Int32}, V::MtlMatrix{T}, k)
+                                                 where {T<:MtlFloat}

Compute the top k values and their corresponding indices column-wise in a matrix A. Return the indices in I and the values in V.

k cannot be greater than 16.

Uses MPSMatrixFindTopK.

See also: topk.

Warn

This interface is experimental, and might change without warning.

source
diff --git a/dev/faq/contributing/index.html b/dev/faq/contributing/index.html index 127045df..bedc1c07 100644 --- a/dev/faq/contributing/index.html +++ b/dev/faq/contributing/index.html @@ -7,4 +7,4 @@ uint i [[thread_position_in_grid]]) { atomic_store_explicit(&out[i], 0.0f, memory_order_relaxed); -}

To compile with Metal's tools and emit human-readable IR, run something roughly along the lines of: xcrun metal -S -emit-llvm dummy_kernel.metal

This will create a .ll file that you can then parse for whatever information you need. Be sure to double-check the metadata at the bottom for any significant changes your functionality introduces.

Test with different types and configurations to see what changes are caused. Also ensure that when writing very simple kernels, whatever you're interested in doesn't get optimized away. Double-check that the kernel's IR makes sense for what you wrote.

Metal Performance Shaders

Metal exposes a special interface to its library of optimized kernels. Rather than accepting the normal set of input GPU data structures, it requires special MPS datatypes that assume row-major memory layout. As this is not the Julia default, adapt accordingly. Adding MPS functionality should be mostly straightforward, so this can be an easy entry point to helping. To get started, you can have a look at the Metal Performance Shaders Documentation from Apple.

Exposing your Interface

There are varying degrees of user-facing interfaces from Metal.jl. At the lowest level is Metal.MTL.xxx. This is for low-level functionality close to or at bare Objective-C, or things that a normal user wouldn't directly be using. Metal.MPS.xxx is for Metal Performance Shader specifics (like MPSMatrix). Next, is Metal.xxx. This is for higher-level, usually pure-Julian functionality (like device()). The only thing beyond this is exporting into the global namespace. That would be useful for uniquely-named functions/structures/macros with clear and common use-cases (MtlArray or @metal).

Additionally, you can override non-Metal.jl functions like LinearAlgebra.mul! seen here. This is essentially (ab)using multiple dispatch to specialize for certain cases (usually for more performant execution).

If your function is only available from within GPU kernels (like thread indexing intrinsics). Be sure to properly annotate with @device_function to ensure that calling from the host doesn't kill your Julia process.

Generally, think about how frequently you expect your addition to be used, how complex its use-case is, and whether or not it clashes/reimplements/optimizes existing functionality from outside Metal.jl. Put it behind the corresponding interface.

Creating Tests

As it's good practice, and JuliaGPU has great CI/CD workflows, your addition should have associated tests to ensure correctness and edge cases. Look to existing examples under the test folder for initial guidance, and be sure to create tests for all valid types. Any new Julia file in this folder will be ran as its own testset. If you feel your tests don't fit in any existing place, you'll probably want to create a new file with an appropriate name.

Running a Subset of the Existing Tests

Sometimes you won't want to run the entire testsuite. You may just want to run the tests for your new functionality. To do that, you can either pass the name of the testset to the test/runtests.jl script: julia --project=test test/runtests.jl metal or you can isolate test files by running them alone after running the test/setup.jl script: julia --project=test -L test/setup.jl test/metal.jl

Thank You and Good Luck

Open-source projects like this only happen because people like you are willing to spend their free time helping out. Most anything you're able to do is helpful, but if you get stuck, seek guidance from Slack or Discourse. Don't feel like your contribution has to be perfect. If you put in effort and make progress, there will likely be some senior developer willing to polish your code before merging. Open-source software is a team effort...welcome to the team!

+}

To compile with Metal's tools and emit human-readable IR, run something roughly along the lines of: xcrun metal -S -emit-llvm dummy_kernel.metal

This will create a .ll file that you can then parse for whatever information you need. Be sure to double-check the metadata at the bottom for any significant changes your functionality introduces.

Test with different types and configurations to see what changes are caused. Also ensure that when writing very simple kernels, whatever you're interested in doesn't get optimized away. Double-check that the kernel's IR makes sense for what you wrote.

Metal Performance Shaders

Metal exposes a special interface to its library of optimized kernels. Rather than accepting the normal set of input GPU data structures, it requires special MPS datatypes that assume row-major memory layout. As this is not the Julia default, adapt accordingly. Adding MPS functionality should be mostly straightforward, so this can be an easy entry point to helping. To get started, you can have a look at the Metal Performance Shaders Documentation from Apple.

Exposing your Interface

There are varying degrees of user-facing interfaces from Metal.jl. At the lowest level is Metal.MTL.xxx. This is for low-level functionality close to or at bare Objective-C, or things that a normal user wouldn't directly be using. Metal.MPS.xxx is for Metal Performance Shader specifics (like MPSMatrix). Next, is Metal.xxx. This is for higher-level, usually pure-Julian functionality (like device()). The only thing beyond this is exporting into the global namespace. That would be useful for uniquely-named functions/structures/macros with clear and common use-cases (MtlArray or @metal).

Additionally, you can override non-Metal.jl functions like LinearAlgebra.mul! seen here. This is essentially (ab)using multiple dispatch to specialize for certain cases (usually for more performant execution).

If your function is only available from within GPU kernels (like thread indexing intrinsics). Be sure to properly annotate with @device_function to ensure that calling from the host doesn't kill your Julia process.

Generally, think about how frequently you expect your addition to be used, how complex its use-case is, and whether or not it clashes/reimplements/optimizes existing functionality from outside Metal.jl. Put it behind the corresponding interface.

Creating Tests

As it's good practice, and JuliaGPU has great CI/CD workflows, your addition should have associated tests to ensure correctness and edge cases. Look to existing examples under the test folder for initial guidance, and be sure to create tests for all valid types. Any new Julia file in this folder will be ran as its own testset. If you feel your tests don't fit in any existing place, you'll probably want to create a new file with an appropriate name.

Running a Subset of the Existing Tests

Sometimes you won't want to run the entire testsuite. You may just want to run the tests for your new functionality. To do that, you can either pass the name of the testset to the test/runtests.jl script: julia --project=test test/runtests.jl metal or you can isolate test files by running them alone after running the test/setup.jl script: julia --project=test -L test/setup.jl test/metal.jl

Thank You and Good Luck

Open-source projects like this only happen because people like you are willing to spend their free time helping out. Most anything you're able to do is helpful, but if you get stuck, seek guidance from Slack or Discourse. Don't feel like your contribution has to be perfect. If you put in effort and make progress, there will likely be some senior developer willing to polish your code before merging. Open-source software is a team effort...welcome to the team!

diff --git a/dev/faq/faq/index.html b/dev/faq/faq/index.html index 63e36408..86449158 100644 --- a/dev/faq/faq/index.html +++ b/dev/faq/faq/index.html @@ -1,2 +1,2 @@ -Frequently Asked Questions · Metal.jl
+Frequently Asked Questions · Metal.jl
diff --git a/dev/index.html b/dev/index.html index 7eb10692..e591deec 100644 --- a/dev/index.html +++ b/dev/index.html @@ -6,4 +6,4 @@ # smoke test using Metal Metal.versioninfo()

If you want to ensure everything works as expected, you can execute the test suite.

using Pkg
-Pkg.test("Metal")

The following resources may also be of interest (although are mainly focused on the CUDA GPU backend):

Contributing

If you want to help improve this package, look at the contributing page for more details.

Acknowledgements

The Julia Metal stack has been a collaborative effort by many individuals. Significant contributions have been made by the following individuals:

Supporting and Citing

Some of the software in this ecosystem was developed as part of academic research. If you would like to help support it, please star the repository as such metrics may help us secure funding in the future. If you use our software as part of your research, teaching, or other activities, we would be grateful if you could cite our work. The CITATION.cff file in the root of this repository lists the relevant papers.

+Pkg.test("Metal")

The following resources may also be of interest (although are mainly focused on the CUDA GPU backend):

Contributing

If you want to help improve this package, look at the contributing page for more details.

Acknowledgements

The Julia Metal stack has been a collaborative effort by many individuals. Significant contributions have been made by the following individuals:

Supporting and Citing

Some of the software in this ecosystem was developed as part of academic research. If you would like to help support it, please star the repository as such metrics may help us secure funding in the future. If you use our software as part of your research, teaching, or other activities, we would be grateful if you could cite our work. The CITATION.cff file in the root of this repository lists the relevant papers.

diff --git a/dev/objects.inv b/dev/objects.inv index 98685638..641165c2 100644 Binary files a/dev/objects.inv and b/dev/objects.inv differ diff --git a/dev/profiling/index.html b/dev/profiling/index.html index 4d7766cd..a00a9906 100644 --- a/dev/profiling/index.html +++ b/dev/profiling/index.html @@ -31,4 +31,4 @@ julia> Metal.@capture @metal threads=length(c) vadd(a, b, c); ... -[ Info: GPU frame capture saved to julia_1.gputrace; open the resulting trace in Xcode +[ Info: GPU frame capture saved to julia_1.gputrace; open the resulting trace in Xcode diff --git a/dev/search_index.js b/dev/search_index.js index 16337237..56c98040 100644 --- a/dev/search_index.js +++ b/dev/search_index.js @@ -1,3 +1,3 @@ var documenterSearchIndex = {"docs": -[{"location":"usage/overview/#UsageOverview","page":"Overview","title":"Overview","text":"","category":"section"},{"location":"usage/overview/","page":"Overview","title":"Overview","text":"The Metal.jl package provides three distinct, but related, interfaces for Metal programming:","category":"page"},{"location":"usage/overview/","page":"Overview","title":"Overview","text":"the MtlArray type: for programming with arrays;\nnative kernel programming capabilities: for writing Metal kernels in Julia;\nMetal API wrappers: for low-level interactions with the Metal libraries.","category":"page"},{"location":"usage/overview/","page":"Overview","title":"Overview","text":"Much of the Julia Metal programming stack can be used by just relying on the MtlArray type, and using platform-agnostic programming patterns like broadcast and other array abstractions. Only once you hit a performance bottleneck, or some missing functionality, you might need to write a custom kernel or use the underlying Metal APIs.","category":"page"},{"location":"usage/overview/#The-MtlArray-type","page":"Overview","title":"The MtlArray type","text":"","category":"section"},{"location":"usage/overview/","page":"Overview","title":"Overview","text":"The MtlArray type is an essential part of the toolchain. Primarily, it is used to manage GPU memory, and copy data from and back to the CPU:","category":"page"},{"location":"usage/overview/","page":"Overview","title":"Overview","text":"a = MtlArray{Int}(undef, 1024)\n\n# essential memory operations, like copying, filling, reshaping, ...\nb = copy(a)\nfill!(b, 0)\n@test b == Metal.zeros(Int, 1024)\n\n# automatic memory management\na = nothing","category":"page"},{"location":"usage/overview/","page":"Overview","title":"Overview","text":"Beyond memory management, there are a whole range of array operations to process your data. This includes several higher-order operations that take other code as arguments, such as map, reduce or broadcast. With these, it is possible to perform kernel-like operations without actually writing your own GPU kernels:","category":"page"},{"location":"usage/overview/","page":"Overview","title":"Overview","text":"a = Metal.zeros(1024)\nb = Metal.ones(1024)\na.^2 .+ sin.(b)","category":"page"},{"location":"usage/array/#Array-programming","page":"Array programming","title":"Array programming","text":"","category":"section"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"DocTestSetup = quote\n using Metal\nend","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"The easiest way to use the GPU's massive parallelism, is by expressing operations in terms of arrays: Metal.jl provides an array type, MtlArray, and many specialized array operations that execute efficiently on the GPU hardware. In this section, we will briefly demonstrate use of the MtlArray type. Since we expose Metal's functionality by implementing existing Julia interfaces on the MtlArray type, you should refer to the upstream Julia documentation for more information on these operations.","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"If you encounter missing functionality, or are running into operations that trigger so-called \"scalar iteration\", have a look at the issue tracker and file a new issue if there's none. Do note that you can always access the underlying Metal APIs by calling into the relevant submodule.","category":"page"},{"location":"usage/array/#Construction-and-Initialization","page":"Array programming","title":"Construction and Initialization","text":"","category":"section"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"The MtlArray type aims to implement the AbstractArray interface, and provide implementations of methods that are commonly used when working with arrays. That means you can construct MtlArrays in the same way as regular Array objects:","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"julia> MtlArray{Int}(undef, 2)\n2-element MtlVector{Int64, Metal.PrivateStorage}:\n 0\n 0\n\njulia> MtlArray{Int}(undef, (1,2))\n1×2 MtlMatrix{Int64, Metal.PrivateStorage}:\n 0 0\n\njulia> similar(ans)\n1×2 MtlMatrix{Int64, Metal.PrivateStorage}:\n 0 0","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"Copying memory to or from the GPU can be expressed using constructors as well, or by calling copyto!:","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"julia> a = MtlArray([1,2])\n2-element MtlVector{Int64, Metal.PrivateStorage}:\n 1\n 2\n\njulia> b = Array(a)\n2-element Vector{Int64}:\n 1\n 2\n\njulia> copyto!(b, a)\n2-element Vector{Int64}:\n 1\n 2","category":"page"},{"location":"usage/array/#Higher-order-abstractions","page":"Array programming","title":"Higher-order abstractions","text":"","category":"section"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"The real power of programming GPUs with arrays comes from Julia's higher-order array abstractions: Operations that take user code as an argument, and specialize execution on it. With these functions, you can often avoid having to write custom kernels. For example, to perform simple element-wise operations you can use map or broadcast:","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"julia> a = MtlArray{Float32}(undef, (1,2));\n\njulia> a .= 5\n1×2 MtlMatrix{Float32, Metal.PrivateStorage}:\n 5.0 5.0\n\njulia> map(sin, a)\n1×2 MtlMatrix{Float32, Metal.PrivateStorage}:\n -0.958924 -0.958924","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"To reduce the dimensionality of arrays, Metal.jl implements the various flavours of (map)reduce(dim):","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"julia> a = Metal.ones(2,3)\n2×3 MtlMatrix{Float32, Metal.PrivateStorage}:\n 1.0 1.0 1.0\n 1.0 1.0 1.0\n\njulia> reduce(+, a)\n6.0f0\n\njulia> mapreduce(sin, *, a; dims=2)\n2×1 MtlMatrix{Float32, Metal.PrivateStorage}:\n 0.59582335\n 0.59582335\n\njulia> b = Metal.zeros(1)\n1-element MtlVector{Float32, Metal.PrivateStorage}:\n 0.0\n\njulia> Base.mapreducedim!(identity, +, b, a)\n1×1 MtlMatrix{Float32, Metal.PrivateStorage}:\n 6.0","category":"page"},{"location":"faq/contributing/#Contributing","page":"Contributing","title":"Contributing","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Metal.jl is an especially accessible GPU backend with the presence of GPUs on Apple's recent popular Macbooks. As a result, an average Julia user can now develop and test GPU-accelerated code locally on their laptop. If you're using this package and see a bug or want some additional functionality, this page is for you. Hopefully this information helps encourage you to contribute to the package yourself.","category":"page"},{"location":"faq/contributing/#What-needs-help?","page":"Contributing","title":"What needs help?","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"If you didn't come to this page with your own feature to add, look at the current issues in the git repo for bugs and requested functionality.","category":"page"},{"location":"faq/contributing/#I'm-a-beginner,-can-I-help?","page":"Contributing","title":"I'm a beginner, can I help?","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Yes, but you may spend more time learning rather than directly contributing at the start. Depending on what your goals are though, this might be desirable. There are differing levels of difficulty when considering contributions to Metal.jl. If you're new to these things, check the issues for \"Good First Issue\" tags, look at the documentation for areas that could be added (beginners are especially good at detecting these sort of deficiencies), or message on the Slack #gpu channel asking for guidance.","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Regardless, if you've never used Metal.jl before, it'd probably be best to gain some exposure to it before trying to contibute. You might run into bugs yourself or discover some area you'd really like to help with.","category":"page"},{"location":"faq/contributing/#General-Workflow-for-Adding-Functionality","page":"Contributing","title":"General Workflow for Adding Functionality","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"If you're adding some functionality that originates from Metal Shading Language (MSL) directly (rather than high-level Julia functionality), the workflow will likely look like the below. If you're adding something that only relies on pure Julia additions, you will skip the first two steps.","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Create low-level, Julia wrappers for the Obj-C interface\nCreate high-level Julia structures and functionality\nCreate tests for added functionality","category":"page"},{"location":"faq/contributing/#Mapping-to-Metal-Intrinsics","page":"Contributing","title":"Mapping to Metal Intrinsics","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Some Metal functions map directly to Apple intermediate representation intrinsics. In this case, wrapping them into Metal.jl is relatively easy. All that needs to be done is to create a mapping from a Julia function via a simple ccall. See the threadgroup barrier implementation for reference.","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"However, the Metal documentation doesn't tell you what the format of the intrinsic names should be. To find this out, you need to create your own test kernel directly in the Metal Shading Language, compile it using Apple's tooling, then view the created intermediate representation (IR).","category":"page"},{"location":"faq/contributing/#Reverse-Engineering-Bare-MSL/Apple-IR","page":"Contributing","title":"Reverse-Engineering Bare MSL/Apple IR","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"First, you need to write an MSL kernel that uses the functionality you're interested in. For example,","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"#include \n\nusing namespace metal;\n\nkernel void dummy_kernel(device volatile atomic_float* out,\n uint i [[thread_position_in_grid]])\n{\n atomic_store_explicit(&out[i], 0.0f, memory_order_relaxed);\n}","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"To compile with Metal's tools and emit human-readable IR, run something roughly along the lines of: xcrun metal -S -emit-llvm dummy_kernel.metal","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"This will create a .ll file that you can then parse for whatever information you need. Be sure to double-check the metadata at the bottom for any significant changes your functionality introduces.","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Test with different types and configurations to see what changes are caused. Also ensure that when writing very simple kernels, whatever you're interested in doesn't get optimized away. Double-check that the kernel's IR makes sense for what you wrote.","category":"page"},{"location":"faq/contributing/#Metal-Performance-Shaders","page":"Contributing","title":"Metal Performance Shaders","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Metal exposes a special interface to its library of optimized kernels. Rather than accepting the normal set of input GPU data structures, it requires special MPS datatypes that assume row-major memory layout. As this is not the Julia default, adapt accordingly. Adding MPS functionality should be mostly straightforward, so this can be an easy entry point to helping. To get started, you can have a look at the Metal Performance Shaders Documentation from Apple.","category":"page"},{"location":"faq/contributing/#Exposing-your-Interface","page":"Contributing","title":"Exposing your Interface","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"There are varying degrees of user-facing interfaces from Metal.jl. At the lowest level is Metal.MTL.xxx. This is for low-level functionality close to or at bare Objective-C, or things that a normal user wouldn't directly be using. Metal.MPS.xxx is for Metal Performance Shader specifics (like MPSMatrix). Next, is Metal.xxx. This is for higher-level, usually pure-Julian functionality (like device()). The only thing beyond this is exporting into the global namespace. That would be useful for uniquely-named functions/structures/macros with clear and common use-cases (MtlArray or @metal).","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Additionally, you can override non-Metal.jl functions like LinearAlgebra.mul! seen here. This is essentially (ab)using multiple dispatch to specialize for certain cases (usually for more performant execution).","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"If your function is only available from within GPU kernels (like thread indexing intrinsics). Be sure to properly annotate with @device_function to ensure that calling from the host doesn't kill your Julia process.","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Generally, think about how frequently you expect your addition to be used, how complex its use-case is, and whether or not it clashes/reimplements/optimizes existing functionality from outside Metal.jl. Put it behind the corresponding interface.","category":"page"},{"location":"faq/contributing/#Creating-Tests","page":"Contributing","title":"Creating Tests","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"As it's good practice, and JuliaGPU has great CI/CD workflows, your addition should have associated tests to ensure correctness and edge cases. Look to existing examples under the test folder for initial guidance, and be sure to create tests for all valid types. Any new Julia file in this folder will be ran as its own testset. If you feel your tests don't fit in any existing place, you'll probably want to create a new file with an appropriate name.","category":"page"},{"location":"faq/contributing/#Running-a-Subset-of-the-Existing-Tests","page":"Contributing","title":"Running a Subset of the Existing Tests","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Sometimes you won't want to run the entire testsuite. You may just want to run the tests for your new functionality. To do that, you can either pass the name of the testset to the test/runtests.jl script: julia --project=test test/runtests.jl metal or you can isolate test files by running them alone after running the test/setup.jl script: julia --project=test -L test/setup.jl test/metal.jl","category":"page"},{"location":"faq/contributing/#Thank-You-and-Good-Luck","page":"Contributing","title":"Thank You and Good Luck","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Open-source projects like this only happen because people like you are willing to spend their free time helping out. Most anything you're able to do is helpful, but if you get stuck, seek guidance from Slack or Discourse. Don't feel like your contribution has to be perfect. If you put in effort and make progress, there will likely be some senior developer willing to polish your code before merging. Open-source software is a team effort...welcome to the team!","category":"page"},{"location":"api/kernel/#Kernel-programming","page":"Kernel programming","title":"Kernel programming","text":"","category":"section"},{"location":"api/kernel/","page":"Kernel programming","title":"Kernel programming","text":"This section lists the package's public functionality that corresponds to special Metal functions for use in device code. For more information about these functions, please consult the Metal Shading Language specification.","category":"page"},{"location":"api/kernel/","page":"Kernel programming","title":"Kernel programming","text":"This is made possible by interfacing with the Metal libraries by wrapping a subset of the ObjectiveC APIs using ObjectiveC.jl. These low-level wrappers are available in the MTL submodule exported by Metal.jl.","category":"page"},{"location":"api/kernel/#Indexing-and-dimensions","page":"Kernel programming","title":"Indexing and dimensions","text":"","category":"section"},{"location":"api/kernel/","page":"Kernel programming","title":"Kernel programming","text":"thread_execution_width\nthread_index_in_quadgroup\nthread_index_in_simdgroup\nthread_index_in_threadgroup\nthread_position_in_grid_1d\nthread_position_in_threadgroup_1d\nthreadgroup_position_in_grid_1d\nthreadgroups_per_grid_1d\nthreads_per_grid_1d\nthreads_per_simdgroup\nthreads_per_threadgroup_1d\nsimdgroups_per_threadgroup\nsimdgroup_index_in_threadgroup\nquadgroup_index_in_threadgroup\nquadgroups_per_threadgroup\ngrid_size_1d\ngrid_origin_1d","category":"page"},{"location":"api/kernel/#Metal.thread_execution_width","page":"Kernel programming","title":"Metal.thread_execution_width","text":"thread_execution_width()::UInt32\n\nReturn the execution width of the compute unit.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.thread_index_in_quadgroup","page":"Kernel programming","title":"Metal.thread_index_in_quadgroup","text":"thread_index_in_quadgroup()::UInt32\n\nReturn the index of the current thread in its quadgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.thread_index_in_simdgroup","page":"Kernel programming","title":"Metal.thread_index_in_simdgroup","text":"thread_index_in_simdgroup()::UInt32\n\nReturn the index of the current thread in its simdgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.thread_index_in_threadgroup","page":"Kernel programming","title":"Metal.thread_index_in_threadgroup","text":"thread_index_in_threadgroup()::UInt32\n\nReturn the index of the current thread in its threadgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.thread_position_in_grid_1d","page":"Kernel programming","title":"Metal.thread_position_in_grid_1d","text":"thread_position_in_grid_1d()::UInt32\nthread_position_in_grid_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}\nthread_position_in_grid_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}\n\nReturn the current thread's position in an N-dimensional grid of threads.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.thread_position_in_threadgroup_1d","page":"Kernel programming","title":"Metal.thread_position_in_threadgroup_1d","text":"thread_position_in_threadgroup_1d()::UInt32\nthread_position_in_threadgroup_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}\nthread_position_in_threadgroup_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}\n\nReturn the current thread's unique position within a threadgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.threadgroup_position_in_grid_1d","page":"Kernel programming","title":"Metal.threadgroup_position_in_grid_1d","text":"threadgroup_position_in_grid_1d()::UInt32\nthreadgroup_position_in_grid_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}\nthreadgroup_position_in_grid_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}\n\nReturn the current threadgroup's unique position within the grid.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.threadgroups_per_grid_1d","page":"Kernel programming","title":"Metal.threadgroups_per_grid_1d","text":"threadgroups_per_grid_1d()::UInt32\nthreadgroups_per_grid_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}\nthreadgroups_per_grid_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}\n\nReturn the number of threadgroups per grid.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.threads_per_grid_1d","page":"Kernel programming","title":"Metal.threads_per_grid_1d","text":"threads_per_grid_1d()::UInt32\nthreads_per_grid_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}\nthreads_per_grid_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}\n\nReturn the grid size.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.threads_per_simdgroup","page":"Kernel programming","title":"Metal.threads_per_simdgroup","text":"threads_per_simdgroup()::UInt32\n\nReturn the thread execution width of a simdgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.threads_per_threadgroup_1d","page":"Kernel programming","title":"Metal.threads_per_threadgroup_1d","text":"threads_per_threadgroup_1d()::UInt32\nthreads_per_threadgroup_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}\nthreads_per_threadgroup_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}\n\nReturn the thread execution width of a threadgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.simdgroups_per_threadgroup","page":"Kernel programming","title":"Metal.simdgroups_per_threadgroup","text":"simdgroups_per_threadgroup()::UInt32\n\nReturn the simdgroup execution width of a threadgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.simdgroup_index_in_threadgroup","page":"Kernel programming","title":"Metal.simdgroup_index_in_threadgroup","text":"simdgroup_index_in_threadgroup()::UInt32\n\nReturn the index of a simdgroup within a threadgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.quadgroup_index_in_threadgroup","page":"Kernel programming","title":"Metal.quadgroup_index_in_threadgroup","text":"quadgroup_index_in_threadgroup()::UInt32\n\nReturn the index of a quadgroup within a threadgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.quadgroups_per_threadgroup","page":"Kernel programming","title":"Metal.quadgroups_per_threadgroup","text":"quadgroups_per_threadgroup()::UInt32\n\nReturn the quadgroup execution width of a threadgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.grid_size_1d","page":"Kernel programming","title":"Metal.grid_size_1d","text":"grid_size_1d()::UInt32\ngrid_size_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}\ngrid_size_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}\n\nReturn maximum size of the grid for threads that read per-thread stage-in data.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.grid_origin_1d","page":"Kernel programming","title":"Metal.grid_origin_1d","text":"grid_origin_1d()::UInt32\ngrid_origin_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}\ngrid_origin_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}\n\nReturn the origin offset of the grid for threads that read per-thread stage-in data.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Device-arrays","page":"Kernel programming","title":"Device arrays","text":"","category":"section"},{"location":"api/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Metal.jl provides a primitive, lightweight array type to manage GPU data organized in an plain, dense fashion. This is the device-counterpart to the MtlArray, and implements (part of) the array interface as well as other functionality for use on the GPU:","category":"page"},{"location":"api/kernel/","page":"Kernel programming","title":"Kernel programming","text":"MtlDeviceArray\nMetal.Const","category":"page"},{"location":"api/kernel/#Metal.MtlDeviceArray","page":"Kernel programming","title":"Metal.MtlDeviceArray","text":"MtlDeviceArray(dims, ptr)\nMtlDeviceArray{T}(dims, ptr)\nMtlDeviceArray{T,A}(dims, ptr)\nMtlDeviceArray{T,A,N}(dims, ptr)\n\nConstruct an N-dimensional dense Metal device array with element type T wrapping a pointer, where N is determined from the length of dims and T is determined from the type of ptr.\n\ndims may be a single scalar, or a tuple of integers corresponding to the lengths in each dimension). If the rank N is supplied explicitly as in Array{T,N}(dims), then it must match the length of dims. The same applies to the element type T, which should match the type of the pointer ptr.\n\n\n\n\n\n","category":"type"},{"location":"api/kernel/#Metal.Const","page":"Kernel programming","title":"Metal.Const","text":"Const(A::MtlDeviceArray)\n\nMark a MtlDeviceArray as constant/read-only and to use the constant address space.\n\nwarning: Warning\nExperimental API. Subject to change without deprecation.\n\n\n\n\n\n","category":"type"},{"location":"api/kernel/#Shared-memory","page":"Kernel programming","title":"Shared memory","text":"","category":"section"},{"location":"api/kernel/","page":"Kernel programming","title":"Kernel programming","text":"MtlThreadGroupArray","category":"page"},{"location":"api/kernel/#Metal.MtlThreadGroupArray","page":"Kernel programming","title":"Metal.MtlThreadGroupArray","text":"MtlThreadGroupArray(::Type{T}, dims)\n\nCreate an array local to each threadgroup launched during kernel execution.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Synchronization","page":"Kernel programming","title":"Synchronization","text":"","category":"section"},{"location":"api/kernel/","page":"Kernel programming","title":"Kernel programming","text":"MemoryFlags\nthreadgroup_barrier\nsimdgroup_barrier","category":"page"},{"location":"api/kernel/#Metal.MemoryFlags","page":"Kernel programming","title":"Metal.MemoryFlags","text":"MemoryFlags\n\nFlags to set the memory synchronization behavior of threadgroup_barrier and simdgroup_barrier.\n\nPossible values:\n\nNone: Set barriers to only act as an execution barrier and not apply a memory fence.\n\nDevice: Ensure the GPU correctly orders the memory operations to device memory\n for threads in the threadgroup or simdgroup.\n\nThreadGroup: Ensure the GPU correctly orders the memory operations to threadgroup\n memory for threads in a threadgroup or simdgroup.\n\nTexture: Ensure the GPU correctly orders the memory operations to texture memory for\n threads in a threadgroup or simdgroup for a texture with the read_write access qualifier.\n\nThreadGroup_ImgBlock: Ensure the GPU correctly orders the memory operations to threadgroup imageblock memory\n for threads in a threadgroup or simdgroup.\n\n\n\n\n\n","category":"type"},{"location":"api/kernel/#Metal.threadgroup_barrier","page":"Kernel programming","title":"Metal.threadgroup_barrier","text":"threadgroup_barrier(flag::MemoryFlags=MemoryFlagNone)\n\nSynchronize all threads in a threadgroup.\n\nPossible flags that affect the memory synchronization behavior are found in MemoryFlags\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.simdgroup_barrier","page":"Kernel programming","title":"Metal.simdgroup_barrier","text":"simdgroup_barrier(flag::MemoryFlags=MemoryFlagNone)\n\nSynchronize all threads in a SIMD-group.\n\nPossible flags that affect the memory synchronization behavior are found in MemoryFlags\n\n\n\n\n\n","category":"function"},{"location":"faq/faq/#Frequently-Asked-Questions","page":"Frequently Asked Questions","title":"Frequently Asked Questions","text":"","category":"section"},{"location":"faq/faq/#Can-you-wrap-this-Metal-API?","page":"Frequently Asked Questions","title":"Can you wrap this Metal API?","text":"","category":"section"},{"location":"faq/faq/","page":"Frequently Asked Questions","title":"Frequently Asked Questions","text":"Most likely. Any help on designing or implementing high-level wrappers for MSL's low-level functionality is greatly appreciated, so please consider contributing your uses of these APIs on the respective repositories.","category":"page"},{"location":"api/mps/#Metal-Performance-Shaders","page":"Metal Performance Shaders","title":"Metal Performance Shaders","text":"","category":"section"},{"location":"api/mps/","page":"Metal Performance Shaders","title":"Metal Performance Shaders","text":"This section lists the package's public functionality that corresponds to the Metal Performance Shaders functions. For more information about these functions, or to see which functions have yet to be implemented in this package, please consult the Metal Performance Shaders Documentation.","category":"page"},{"location":"api/mps/#Matrices-and-Vectors","page":"Metal Performance Shaders","title":"Matrices and Vectors","text":"","category":"section"},{"location":"api/mps/","page":"Metal Performance Shaders","title":"Metal Performance Shaders","text":"MPS.MPSMatrix\nMPS.MPSVector","category":"page"},{"location":"api/mps/#Metal.MPS.MPSMatrix","page":"Metal Performance Shaders","title":"Metal.MPS.MPSMatrix","text":"MPSMatrix(mat::MtlMatrix)\n\nMetal matrix representation used in Performance Shaders.\n\nNote that this results in a transposed view of the input, as Metal stores matrices row-major instead of column-major.\n\n\n\n\n\nMPSMatrix(vec::MtlVector)\n\nMetal matrix representation used in Performance Shaders.\n\nNote that this results in a transposed view of the input, as Metal stores matrices row-major instead of column-major.\n\n\n\n\n\nMPSMatrix(arr::MtlArray{T,3})\n\nMetal batched matrix representation used in Performance Shaders.\n\nNote that this results in a transposed view of the input, as Metal stores matrices row-major instead of column-major.\n\n\n\n\n\n","category":"type"},{"location":"api/mps/#Metal.MPS.MPSVector","page":"Metal Performance Shaders","title":"Metal.MPS.MPSVector","text":"MPSVector(arr::MtlVector)\n\nMetal vector representation used in Performance Shaders.\n\n\n\n\n\n","category":"type"},{"location":"api/mps/#Matrix-Arithmetic-Operators","page":"Metal Performance Shaders","title":"Matrix Arithmetic Operators","text":"","category":"section"},{"location":"api/mps/","page":"Metal Performance Shaders","title":"Metal Performance Shaders","text":"MPS.matmul!\nMPS.matvecmul!\nMPS.topk\nMPS.topk!","category":"page"},{"location":"api/mps/#Metal.MPS.matmul!","page":"Metal Performance Shaders","title":"Metal.MPS.matmul!","text":"matMulMPS(a::MtlMatrix, b::MtlMatrix, c::MtlMatrix, alpha=1, beta=1,\n transpose_left=false, transpose_right=false)\n\nA MPSMatrixMultiplication kernel thay computes: c = alpha * op(a) * beta * op(b) + beta * C\n\nThis function should not typically be used. Rather, use the normal LinearAlgebra interface with any MtlArray and it should be accelerated using Metal Performance Shaders.\n\n\n\n\n\n","category":"function"},{"location":"api/mps/#Metal.MPS.matvecmul!","page":"Metal Performance Shaders","title":"Metal.MPS.matvecmul!","text":"matvecmul!(c::MtlVector, a::MtlMatrix, b::MtlVector, alpha=1, beta=1, transpose=false)\n\nA MPSMatrixVectorMultiplication kernel thay computes: c = alpha * op(a) * b + beta * c\n\nThis function should not typically be used. Rather, use the normal LinearAlgebra interface with any MtlArray and it should be accelerated using Metal Performance Shaders.\n\n\n\n\n\n","category":"function"},{"location":"api/mps/#Metal.MPS.topk","page":"Metal Performance Shaders","title":"Metal.MPS.topk","text":"MPS.topk(A::MtlMatrix{T}, k) where {T<:MtlFloat}\n\nCompute the top k values and their corresponding indices column-wise in a matrix A. Return the indices in I and the values in V.\n\nk cannot be greater than 16.\n\nUses MPSMatrixFindTopK.\n\nSee also: topk!.\n\nwarn: Warn\nThis interface is experimental, and might change without warning.\n\n\n\n\n\n","category":"function"},{"location":"api/mps/#Metal.MPS.topk!","page":"Metal Performance Shaders","title":"Metal.MPS.topk!","text":"MPS.topk!(A::MtlMatrix{T}, I::MtlMatrix{Int32}, V::MtlMatrix{T}, k)\n where {T<:MtlFloat}\n\nCompute the top k values and their corresponding indices column-wise in a matrix A. Return the indices in I and the values in V.\n\nk cannot be greater than 16.\n\nUses MPSMatrixFindTopK.\n\nSee also: topk.\n\nwarn: Warn\nThis interface is experimental, and might change without warning.\n\n\n\n\n\n","category":"function"},{"location":"api/array/#Array-programming","page":"Array programming","title":"Array programming","text":"","category":"section"},{"location":"api/array/","page":"Array programming","title":"Array programming","text":"The Metal array type, MtlArray, generally implements the Base array interface and all of its expected methods.","category":"page"},{"location":"api/array/","page":"Array programming","title":"Array programming","text":"However, there is the special function mtl for transferring an array over to the gpu. For compatibility reasons, it will automatically convert arrays of Float64 to Float32.","category":"page"},{"location":"api/array/","page":"Array programming","title":"Array programming","text":"mtl\nMtlArray\nMtlVector\nMtlMatrix\nMtlVecOrMat","category":"page"},{"location":"api/array/#Metal.mtl","page":"Array programming","title":"Metal.mtl","text":"mtl(A; storage=Metal.PrivateStorage)\n\nstorage can be Metal.PrivateStorage (default), Metal.SharedStorage, or Metal.ManagedStorage.\n\nOpinionated GPU array adaptor, which may alter the element type T of arrays:\n\nFor T<:AbstractFloat, it makes a MtlArray{Float32} for performance and compatibility reasons (except for Float16).\nFor T<:Complex{<:AbstractFloat} it makes a MtlArray{ComplexF32}.\nFor other isbitstype(T), it makes a MtlArray{T}.\n\nBy contrast, MtlArray(A) never changes the element type.\n\nUses Adapt.jl to act inside some wrapper structs.\n\nExamples\n\njulia> mtl(ones(3)')\n1×3 adjoint(::MtlVector{Float32, Metal.PrivateStorage}) with eltype Float32:\n 1.0 1.0 1.0\n\njulia> mtl(zeros(1,3); storage=Metal.SharedStorage)\n1×3 MtlMatrix{Float32, Metal.SharedStorage}:\n 0.0 0.0 0.0\n\njulia> mtl(1:3)\n1:3\n\njulia> MtlArray(1:3)\n3-element MtlVector{Int64, Metal.PrivateStorage}:\n 1\n 2\n 3\n\n\n\n\n\n","category":"function"},{"location":"api/array/#Metal.MtlArray","page":"Array programming","title":"Metal.MtlArray","text":"MtlArray{T,N,S} <: AbstractGPUArray{T,N}\n\nN-dimensional Metal array with storage mode S and elements of type T.\n\nS can be Metal.PrivateStorage (default), Metal.SharedStorage, or Metal.ManagedStorage.\n\nSee the Array Programming section of the Metal.jl docs for more details.\n\n\n\n\n\n","category":"type"},{"location":"api/array/#Metal.MtlVector","page":"Array programming","title":"Metal.MtlVector","text":"MtlVector{T,S} <: AbstractGPUVector{T}\n\nOne-dimensional array with elements of type T for use with Apple Metal-compatible GPUs. Alias for MtlArray{T,1,S}.\n\nSee also Vector(@ref), and the Array Programming section of the Metal.jl docs for more details.\n\n\n\n\n\n","category":"type"},{"location":"api/array/#Metal.MtlMatrix","page":"Array programming","title":"Metal.MtlMatrix","text":"MtlMatrix{T,S} <: AbstractGPUMatrix{T}\n\nTwo-dimensional array with elements of type T for use with Apple Metal-compatible GPUs. Alias for MtlArray{T,2,S}.\n\nSee also Matrix(@ref), and the Array Programming section of the Metal.jl docs for more details.\n\n\n\n\n\n","category":"type"},{"location":"api/array/#Metal.MtlVecOrMat","page":"Array programming","title":"Metal.MtlVecOrMat","text":"MtlVecOrMat{T,S}\n\nUnion type of MtlVector{T,S} and MtlMatrix{T,S} which allows functions to accept either an MtlMatrix or an MtlVector.\n\nSee also VecOrMat(@ref) for examples.\n\n\n\n\n\n","category":"type"},{"location":"api/array/#Storage-modes","page":"Array programming","title":"Storage modes","text":"","category":"section"},{"location":"api/array/","page":"Array programming","title":"Array programming","text":"The Metal API has various storage modes that dictate how a resource can be accessed. MtlArrays are Metal.PrivateStorage by default, but they can also be Metal.SharedStorage or Metal.ManagedStorage. For more information on storage modes, see the official Metal documentation.","category":"page"},{"location":"api/array/","page":"Array programming","title":"Array programming","text":"Metal.PrivateStorage\nMetal.SharedStorage\nMetal.ManagedStorage","category":"page"},{"location":"api/array/#Metal.MTL.PrivateStorage","page":"Array programming","title":"Metal.MTL.PrivateStorage","text":"struct Metal.PrivateStorage <: MTL.StorageMode\n\nUsed to indicate that the resource is stored using MTLStorageModePrivate in memory.\n\nFor more information on Metal storage modes, refer to the official Metal documentation.\n\nSee also Metal.SharedStorage and Metal.ManagedStorage.\n\n\n\n\n\n","category":"type"},{"location":"api/array/#Metal.MTL.SharedStorage","page":"Array programming","title":"Metal.MTL.SharedStorage","text":"struct Metal.SharedStorage <: MTL.StorageMode\n\nUsed to indicate that the resource is stored using MTLStorageModeShared in memory.\n\nFor more information on Metal storage modes, refer to the official Metal documentation.\n\nSee also Metal.PrivateStorage and Metal.ManagedStorage.\n\n\n\n\n\n","category":"type"},{"location":"api/array/#Metal.MTL.ManagedStorage","page":"Array programming","title":"Metal.MTL.ManagedStorage","text":"struct Metal.ManagedStorage <: MTL.StorageMode\n\nUsed to indicate that the resource is stored using MTLStorageModeManaged in memory.\n\nFor more information on Metal storage modes, refer to the official Metal documentation.\n\nSee also Metal.SharedStorage and Metal.PrivateStorage.\n\n\n\n\n\n","category":"type"},{"location":"api/array/","page":"Array programming","title":"Array programming","text":"There also exist the following convenience functions to check if an MtlArray is using a specific storage mode:","category":"page"},{"location":"api/array/","page":"Array programming","title":"Array programming","text":"is_private\nis_shared\nis_managed","category":"page"},{"location":"api/array/#Metal.is_private","page":"Array programming","title":"Metal.is_private","text":"is_private(A::MtlArray) -> Bool\n\nReturns true if A has storage mode Metal.PrivateStorage.\n\nSee also is_shared and is_managed.\n\n\n\n\n\n","category":"function"},{"location":"api/array/#Metal.is_shared","page":"Array programming","title":"Metal.is_shared","text":"is_shared(A::MtlArray) -> Bool\n\nReturns true if A has storage mode Metal.SharedStorage.\n\nSee also is_private and is_managed.\n\n\n\n\n\n","category":"function"},{"location":"api/array/#Metal.is_managed","page":"Array programming","title":"Metal.is_managed","text":"is_managed(A::MtlArray) -> Bool\n\nReturns true if A has storage mode Metal.ManagedStorage.\n\nSee also is_shared and is_private.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Essentials","page":"Essentials","title":"Essentials","text":"","category":"section"},{"location":"api/essentials/#Versions-and-Support","page":"Essentials","title":"Versions and Support","text":"","category":"section"},{"location":"api/essentials/","page":"Essentials","title":"Essentials","text":"Metal.macos_version\nMetal.darwin_version\nMetal.metal_support\nMetal.metallib_support\nMetal.air_support","category":"page"},{"location":"api/essentials/#Metal.macos_version","page":"Essentials","title":"Metal.macos_version","text":"Metal.macos_version() -> VersionNumber\n\nReturns the host macOS version.\n\nSee also Metal.darwin_version.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.darwin_version","page":"Essentials","title":"Metal.darwin_version","text":"Metal.darwin_version() -> VersionNumber\n\nReturns the host Darwin kernel version.\n\nSee also Metal.macos_version.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.metal_support","page":"Essentials","title":"Metal.metal_support","text":"Metal.metal_support() -> VersionNumber\n\nReturns the highest supported version for the Metal Shading Language.\n\nSee also Metal.metallib_support and Metal.air_support.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.metallib_support","page":"Essentials","title":"Metal.metallib_support","text":"Metal.metallib_support() -> VersionNumber\n\nReturns the highest supported version for the metallib file format.\n\nSee also Metal.air_support and Metal.metal_support.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.air_support","page":"Essentials","title":"Metal.air_support","text":"Metal.air_support() -> VersionNumber\n\nReturns the highest supported version for the embedded AIR bitcode format.\n\nSee also Metal.metallib_support and Metal.metal_support.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Global-State","page":"Essentials","title":"Global State","text":"","category":"section"},{"location":"api/essentials/","page":"Essentials","title":"Essentials","text":"Metal.device!\nMetal.devices\nMetal.device\nMetal.global_queue\nMetal.synchronize\nMetal.device_synchronize","category":"page"},{"location":"api/essentials/#Metal.device!","page":"Essentials","title":"Metal.device!","text":"device!(dev::MTLDevice)\n\nSets the Metal GPU device associated with the current Julia task.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.MTL.devices","page":"Essentials","title":"Metal.MTL.devices","text":"devices()\n\nGet an iterator for the compute devices.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.device","page":"Essentials","title":"Metal.device","text":"device()::MTLDevice\n\nReturn the Metal GPU device associated with the current Julia task.\n\nSince all M-series systems currently only externally show a single GPU, this function effectively returns the only system GPU.\n\n\n\n\n\ndevice(<:MtlArray)\n\nGet the Metal device for an MtlArray.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.global_queue","page":"Essentials","title":"Metal.global_queue","text":"global_queue(dev::MTLDevice)::MTLCommandQueue\n\nReturn the Metal command queue associated with the current Julia thread.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.synchronize","page":"Essentials","title":"Metal.synchronize","text":"synchronize(queue)\n\nWait for currently committed GPU work on this queue to finish.\n\nCreate a new MTLCommandBuffer from the global command queue, commit it to the queue, and simply wait for it to be completed. Since command buffers should execute in a First-In-First-Out manner, this synchronizes the GPU.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.device_synchronize","page":"Essentials","title":"Metal.device_synchronize","text":"device_synchronize()\n\nSynchronize all committed GPU work across all global queues\n\n\n\n\n\n","category":"function"},{"location":"usage/kernel/#Kernel-programming","page":"Kernel programming","title":"Kernel programming","text":"","category":"section"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Metal.jl is based off of Apple's Metal Shading Language (MSL) and Metal framework. The interface allows you to utilize the graphics and computing power of Mac GPUs. Like many other GPU frameworks, its history is rooted in graphics processing but has found use in computing/general purpose GPU (GPGPU) applications.","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"The most fundamental idea of programming GPUs (when compared to serial CPU programming) is its parallelism. A GPU function (kernel), when called, is not just ran once in isolation. Rather, numerous (often thousands to millions) psuedo-independent instances (called threads) of the kernel are executed in parallel. These threads are arranged in a hierarchy that allows for varying levels of synchronization. For Metal, the hierarchy is as follows:","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Thread: A single execution unit of the kernel\nThreadgroup: A collection of threads that share a common block of memory and synchronization","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"barriers","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Grid: A collection of threadgroups","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"The threadgroup and grid sizes are set by the user when launching the GPU kernel. There are upper limits determined by the targeted hardware, and the sizes can be 1, 2, or 3-dimensional. For Metal.jl, these sizes are set using the @metal macro's keyword arguments. The grid keyword determines the grid size while the threads keyword determines the threadgroup size.","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"For example, given a 10x10x3 image that you want to run a function independently on each pixel, the kernel launch code might look like the following: @metal threads=(10,10) groups=3 my_kernel(gpu_image_array) This would launch 3 separate threadgroups of 100 threads each (10 in the first dimension and 10 in the second dimension)","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"There is also additional hierarchy layers that consists of small groups of threads that execute in lockstep called waves/SIMD groups/wavefronts* and quadgroups. However, the basic three-tier hierarchy is enough to get started.","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Here is a helpful link with good visualizations of Metal's thread hierarchy (also covering SIMD groups).","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Each thread has its own set of private variables. Most importantly, each thread has associated unique indices to identify itself within its threadgroup and grid. These are traditionally what are used to differentiate execution across threads. You can also query what the grid and threadgroup sizes are as well.","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"For Metal.jl, these values are accessed via the following functions:","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"thread_index_in_threadgroup()\ngrid_size_Xd()\nthread_position_in_grid_Xd()\nthread_position_in_threadgroup_Xd()\nthreadgroup_position_in_grid_Xd()\nthreadgroups_per_grid_Xd()\nthreads_per_grid_Xd()\nthreads_per_threadgroup_Xd()","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Where 'X' is 1, 2, or 3 according to the number of dimensions requested.","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Using these in a kernel (taken directly from the vadd example):","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"function vadd(a, b, c)\n i = thread_position_in_grid_1d()\n c[i] = a[i] + b[i]\n return\nend","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"This kernel takes in three vectors (a,b,c) all of the same length and stores the element-wise sum of a and b into c. Each thread in this kernel gets its unique position in the grid (arrangement of all threadgroups) and stores this value into the variable i which is then used as the index into the vectors. Thus, each thread is computing one sum and storing the result in the output vector.","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"To ensure this kernel functions properly, we have to launch it with exactly as many threads as the length of the vectors. If we under or over-launch threads, the result could be incorrect.","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"An example of a good launch:","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"len = prod(size(d_a))\n@metal threads=len vadd(d_a, d_b, d_c)","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Additional notes:","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Kernels must always return nothing\nKernels are asynchronous. To synchronize, use the Metal.@sync macro.","category":"page"},{"location":"usage/kernel/#Other-Helpful-Links","page":"Kernel programming","title":"Other Helpful Links","text":"","category":"section"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Metal Shading Language Specification An Introduction to GPU Programming course from University of Illinois (primarily in CUDA, but the concepts are transferable)","category":"page"},{"location":"api/compiler/#Compiler","page":"Compiler","title":"Compiler","text":"","category":"section"},{"location":"api/compiler/#Execution","page":"Compiler","title":"Execution","text":"","category":"section"},{"location":"api/compiler/","page":"Compiler","title":"Compiler","text":"The main entry-point to the compiler is the @metal macro:","category":"page"},{"location":"api/compiler/","page":"Compiler","title":"Compiler","text":"@metal","category":"page"},{"location":"api/compiler/#Metal.@metal","page":"Compiler","title":"Metal.@metal","text":"@metal threads=... groups=... [kwargs...] func(args...)\n\nHigh-level interface for executing code on a GPU.\n\nThe @metal macro should prefix a call, with func a callable function or object that should return nothing. It will be compiled to a Metal function upon first use, and to a certain extent arguments will be converted and managed automatically using mtlconvert. Finally, a call to mtlcall is performed, creating a command buffer in the current global command queue then committing it.\n\nThere is one supported keyword argument that influences the behavior of @metal:\n\nlaunch: whether to launch this kernel, defaults to true. If false the returned kernel object should be launched by calling it and passing arguments again.\nname: the name of the kernel in the generated code. Defaults to an automatically- generated name.\nqueue: the command queue to use for this kernel. Defaults to the global command queue.\n\n\n\n\n\n","category":"macro"},{"location":"api/compiler/","page":"Compiler","title":"Compiler","text":"If needed, you can use a lower-level API that lets you inspect the compiler kernel:","category":"page"},{"location":"api/compiler/","page":"Compiler","title":"Compiler","text":"Metal.mtlconvert\nMetal.mtlfunction","category":"page"},{"location":"api/compiler/#Metal.mtlconvert","page":"Compiler","title":"Metal.mtlconvert","text":"mtlconvert(x, [cce])\n\nThis function is called for every argument to be passed to a kernel, allowing it to be converted to a GPU-friendly format. By default, the function does nothing and returns the input object x as-is.\n\nDo not add methods to this function, but instead extend the underlying Adapt.jl package and register methods for the the Metal.Adaptor type.\n\n\n\n\n\n","category":"function"},{"location":"api/compiler/#Metal.mtlfunction","page":"Compiler","title":"Metal.mtlfunction","text":"mtlfunction(f, tt=Tuple{}; kwargs...)\n\nLow-level interface to compile a function invocation for the currently-active GPU, returning a callable kernel object. For a higher-level interface, use @metal.\n\nThe output of this function is automatically cached, i.e. you can simply call mtlfunction in a hot path without degrading performance. New code will be generated automatically when the function changes, or when different types or keyword arguments are provided.\n\n\n\n\n\n","category":"function"},{"location":"api/compiler/#Reflection","page":"Compiler","title":"Reflection","text":"","category":"section"},{"location":"api/compiler/","page":"Compiler","title":"Compiler","text":"If you want to inspect generated code, you can use macros that resemble functionality from the InteractiveUtils standard library:","category":"page"},{"location":"api/compiler/","page":"Compiler","title":"Compiler","text":"@device_code_lowered\n@device_code_typed\n@device_code_warntype\n@device_code_llvm\n@device_code_native\n@device_code_agx\n@device_code","category":"page"},{"location":"api/compiler/","page":"Compiler","title":"Compiler","text":"For more information, please consult the GPUCompiler.jl documentation. code_agx is actually code_native:","category":"page"},{"location":"#MacOS-GPU-programming-in-Julia","page":"Home","title":"MacOS GPU programming in Julia","text":"","category":"section"},{"location":"","page":"Home","title":"Home","text":"The Metal.jl package is the main entry point for GPU programming on MacOS in Julia. The package makes it possible to do so at various abstraction levels, from easy-to-use arrays down to hand-written kernels using low-level Metal APIs.","category":"page"},{"location":"","page":"Home","title":"Home","text":"If you have any questions, please feel free to use the #gpu channel on the Julia slack, or the GPU domain of the Julia Discourse.","category":"page"},{"location":"","page":"Home","title":"Home","text":"As this package is still under development, if you spot a bug, please file an issue.","category":"page"},{"location":"#Quick-Start","page":"Home","title":"Quick Start","text":"","category":"section"},{"location":"","page":"Home","title":"Home","text":"Metal.jl ties into your system's existing Metal Shading Language compiler toolchain, so no additional installs are required (unless you want to view profiled GPU operations)","category":"page"},{"location":"","page":"Home","title":"Home","text":"# install the package\nusing Pkg\nPkg.add(\"Metal\")\n\n# smoke test\nusing Metal\nMetal.versioninfo()","category":"page"},{"location":"","page":"Home","title":"Home","text":"If you want to ensure everything works as expected, you can execute the test suite.","category":"page"},{"location":"","page":"Home","title":"Home","text":"using Pkg\nPkg.test(\"Metal\")","category":"page"},{"location":"","page":"Home","title":"Home","text":"The following resources may also be of interest (although are mainly focused on the CUDA GPU backend):","category":"page"},{"location":"","page":"Home","title":"Home","text":"Effectively using GPUs with Julia: slides\nHow Julia is compiled to GPUs: video","category":"page"},{"location":"#Contributing","page":"Home","title":"Contributing","text":"","category":"section"},{"location":"","page":"Home","title":"Home","text":"If you want to help improve this package, look at the contributing page for more details.","category":"page"},{"location":"#Acknowledgements","page":"Home","title":"Acknowledgements","text":"","category":"section"},{"location":"","page":"Home","title":"Home","text":"The Julia Metal stack has been a collaborative effort by many individuals. Significant contributions have been made by the following individuals:","category":"page"},{"location":"","page":"Home","title":"Home","text":"Tim Besard (@maleadt) (lead developer)\nFilippo Vicentini (@PhilipVinc)\nMax Hawkins (@max-Hawkins)","category":"page"},{"location":"#Supporting-and-Citing","page":"Home","title":"Supporting and Citing","text":"","category":"section"},{"location":"","page":"Home","title":"Home","text":"Some of the software in this ecosystem was developed as part of academic research. If you would like to help support it, please star the repository as such metrics may help us secure funding in the future. If you use our software as part of your research, teaching, or other activities, we would be grateful if you could cite our work. The CITATION.cff file in the root of this repository lists the relevant papers.","category":"page"},{"location":"profiling/#Profiling","page":"Profiling","title":"Profiling","text":"","category":"section"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"Profiling GPU code is harder than profiling Julia code executing on the CPU. For one, kernels typically execute asynchronously, and thus require appropriate synchronization when measuring their execution time. Furthermore, because the code executes on a different processor, it is much harder to know what is currently executing.","category":"page"},{"location":"profiling/#Time-measurements","page":"Profiling","title":"Time measurements","text":"","category":"section"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"For robust measurements, it is advised to use the BenchmarkTools.jl package which goes to great lengths to perform accurate measurements. Due to the asynchronous nature of GPUs, you need to ensure the GPU is synchronized at the end of every sample, e.g. by calling synchronize() or, even better, wrapping your code in Metal.@sync:","category":"page"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"Note that the allocations as reported by BenchmarkTools are CPU allocations.","category":"page"},{"location":"profiling/#Application-tracing","page":"Profiling","title":"Application tracing","text":"","category":"section"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"For profiling large applications, simple timings are insufficient. Instead, we want an overview of how and when the GPU was active to avoid times where the device was idle and/or find which kernels needs optimization.","category":"page"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"As we cannot use the Julia profiler for this task, we will use Metal's GPU profiler directly. Use the Metal.@profile macro to surround the code code of interest. This macro tells your system to track GPU calls and usage statistics and will save this information in a temporary folder ending in '.trace'. For later viewing, copy this folder to a stable location or use the 'dir' argument of the profile macro to store the gputrace to a different location directly.","category":"page"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"The resulting trace can be opened with the Instruments app, part of Xcode.","category":"page"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"julia> using Metal\n\njulia> function vadd(a, b, c)\n i = thread_position_in_grid_1d()\n c[i] = a[i] + b[i]\n return\n end\njulia> a = MtlArray([1]); b = MtlArray([2]); c = similar(a);\n\njulia> Metal.@profile @metal threads=length(c) vadd(a, b, c);\n...\n[ Info: System trace saved to julia_3.trace; open the resulting trace in Instruments","category":"page"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"It is possible to augment the trace with additional information by using signposts: Similar to NVTX markers and ranges in CUDA.jl, signpost intervals and events can be used to add respectively time intervals and points of interest to the trace. This can be done by using the signpost functionality from ObjectiveC.jl:","category":"page"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"using ObjectiveC, .OS\n\n@signpost_interval \"My Interval\" begin\n # code to profile\n @signpost_event \"My Event\"\nend","category":"page"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"For more information, e.g. how to pass additional messages to the signposts, or how to use a custom logger, consult the ObjectiveC.jl documentation, or the docstrings of the @signpost_interval and @signpost_event macros.","category":"page"},{"location":"profiling/#Frame-capture","page":"Profiling","title":"Frame capture","text":"","category":"section"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"For more details on specific operations, you can use Metal's frame capture feature to generate a more detailed, and replayable trace of the GPU operations. This requires that Julia is started with the METAL_CAPTURE_ENABLED environment variable set to 1. Frames are captured by wrapping the code of interest in Metal.@capture, and the resulting trace can be opened with Xcode.","category":"page"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"$ METAL_CAPTURE_ENABLED=1 julia\n...\n\njulia> using Metal\n\njulia> function vadd(a, b, c)\n i = thread_position_in_grid_1d()\n c[i] = a[i] + b[i]\n return\n end\n\njulia> a = MtlArray([1]); b = MtlArray([2]); c = similar(a);\n... Metal GPU Frame Capture Enabled\n\njulia> Metal.@capture @metal threads=length(c) vadd(a, b, c);\n...\n[ Info: GPU frame capture saved to julia_1.gputrace; open the resulting trace in Xcode","category":"page"}] +[{"location":"usage/overview/#UsageOverview","page":"Overview","title":"Overview","text":"","category":"section"},{"location":"usage/overview/","page":"Overview","title":"Overview","text":"The Metal.jl package provides three distinct, but related, interfaces for Metal programming:","category":"page"},{"location":"usage/overview/","page":"Overview","title":"Overview","text":"the MtlArray type: for programming with arrays;\nnative kernel programming capabilities: for writing Metal kernels in Julia;\nMetal API wrappers: for low-level interactions with the Metal libraries.","category":"page"},{"location":"usage/overview/","page":"Overview","title":"Overview","text":"Much of the Julia Metal programming stack can be used by just relying on the MtlArray type, and using platform-agnostic programming patterns like broadcast and other array abstractions. Only once you hit a performance bottleneck, or some missing functionality, you might need to write a custom kernel or use the underlying Metal APIs.","category":"page"},{"location":"usage/overview/#The-MtlArray-type","page":"Overview","title":"The MtlArray type","text":"","category":"section"},{"location":"usage/overview/","page":"Overview","title":"Overview","text":"The MtlArray type is an essential part of the toolchain. Primarily, it is used to manage GPU memory, and copy data from and back to the CPU:","category":"page"},{"location":"usage/overview/","page":"Overview","title":"Overview","text":"a = MtlArray{Int}(undef, 1024)\n\n# essential memory operations, like copying, filling, reshaping, ...\nb = copy(a)\nfill!(b, 0)\n@test b == Metal.zeros(Int, 1024)\n\n# automatic memory management\na = nothing","category":"page"},{"location":"usage/overview/","page":"Overview","title":"Overview","text":"Beyond memory management, there are a whole range of array operations to process your data. This includes several higher-order operations that take other code as arguments, such as map, reduce or broadcast. With these, it is possible to perform kernel-like operations without actually writing your own GPU kernels:","category":"page"},{"location":"usage/overview/","page":"Overview","title":"Overview","text":"a = Metal.zeros(1024)\nb = Metal.ones(1024)\na.^2 .+ sin.(b)","category":"page"},{"location":"usage/array/#Array-programming","page":"Array programming","title":"Array programming","text":"","category":"section"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"DocTestSetup = quote\n using Metal\n using GPUArrays\n\n import Random\n Random.seed!(1)\n\n Metal.seed!(1)\nend","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"The easiest way to use the GPU's massive parallelism, is by expressing operations in terms of arrays: Metal.jl provides an array type, MtlArray, and many specialized array operations that execute efficiently on the GPU hardware. In this section, we will briefly demonstrate use of the MtlArray type. Since we expose Metal's functionality by implementing existing Julia interfaces on the MtlArray type, you should refer to the upstream Julia documentation for more information on these operations.","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"If you encounter missing functionality, or are running into operations that trigger so-called \"scalar iteration\", have a look at the issue tracker and file a new issue if there's none. Do note that you can always access the underlying Metal APIs by calling into the relevant submodule.","category":"page"},{"location":"usage/array/#Construction-and-Initialization","page":"Array programming","title":"Construction and Initialization","text":"","category":"section"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"The MtlArray type aims to implement the AbstractArray interface, and provide implementations of methods that are commonly used when working with arrays. That means you can construct MtlArrays in the same way as regular Array objects:","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"julia> MtlArray{Int}(undef, 2)\n2-element MtlVector{Int64, Metal.PrivateStorage}:\n 0\n 0\n\njulia> MtlArray{Int}(undef, (1,2))\n1×2 MtlMatrix{Int64, Metal.PrivateStorage}:\n 0 0\n\njulia> similar(ans)\n1×2 MtlMatrix{Int64, Metal.PrivateStorage}:\n 0 0","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"Copying memory to or from the GPU can be expressed using constructors as well, or by calling copyto!:","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"julia> a = MtlArray([1,2])\n2-element MtlVector{Int64, Metal.PrivateStorage}:\n 1\n 2\n\njulia> b = Array(a)\n2-element Vector{Int64}:\n 1\n 2\n\njulia> copyto!(b, a)\n2-element Vector{Int64}:\n 1\n 2","category":"page"},{"location":"usage/array/#Higher-order-abstractions","page":"Array programming","title":"Higher-order abstractions","text":"","category":"section"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"The real power of programming GPUs with arrays comes from Julia's higher-order array abstractions: Operations that take user code as an argument, and specialize execution on it. With these functions, you can often avoid having to write custom kernels. For example, to perform simple element-wise operations you can use map or broadcast:","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"julia> a = MtlArray{Float32}(undef, (1,2));\n\njulia> a .= 5\n1×2 MtlMatrix{Float32, Metal.PrivateStorage}:\n 5.0 5.0\n\njulia> map(sin, a)\n1×2 MtlMatrix{Float32, Metal.PrivateStorage}:\n -0.958924 -0.958924","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"To reduce the dimensionality of arrays, Metal.jl implements the various flavours of (map)reduce(dim):","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"julia> a = Metal.ones(2,3)\n2×3 MtlMatrix{Float32, Metal.PrivateStorage}:\n 1.0 1.0 1.0\n 1.0 1.0 1.0\n\njulia> reduce(+, a)\n6.0f0\n\njulia> mapreduce(sin, *, a; dims=2)\n2×1 MtlMatrix{Float32, Metal.PrivateStorage}:\n 0.59582335\n 0.59582335\n\njulia> b = Metal.zeros(1)\n1-element MtlVector{Float32, Metal.PrivateStorage}:\n 0.0\n\njulia> Base.mapreducedim!(identity, +, b, a)\n1×1 MtlMatrix{Float32, Metal.PrivateStorage}:\n 6.0","category":"page"},{"location":"usage/array/#Random-numbers","page":"Array programming","title":"Random numbers","text":"","category":"section"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"Base's convenience functions for generating random numbers are available in Metal as well:","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"julia> Metal.rand(2)\n2-element MtlVector{Float32, Metal.PrivateStorage}:\n 0.89025915\n 0.8946847\n\njulia> Metal.randn(Float32, 2, 1)\n2×1 MtlMatrix{Float32, Metal.PrivateStorage}:\n 1.2279074\n 1.2518331","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"Behind the scenes, these random numbers come from two different generators: one backed by Metal Performance Shaders, another by using the GPUArrays.jl random methods. Operations on these generators are implemented using methods from the Random standard library:","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"julia> using Random, GPUArrays\n\njulia> a = Random.rand(MPS.default_rng(), Float32, 1)\n1-element MtlVector{Float32, Metal.PrivateStorage}:\n 0.89025915\n\njulia> a = Random.rand!(GPUArrays.default_rng(MtlArray), a)\n1-element MtlVector{Float32, Metal.PrivateStorage}:\n 0.0705002","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"note: Note\nMPSMatrixRandom functionality requires Metal.jl >= v1.4","category":"page"},{"location":"usage/array/","page":"Array programming","title":"Array programming","text":"warning: Warning\nRandom.rand!(::MPS.RNG, args...) and Random.randn!(::MPS.RNG, args...) have a framework limitation that requires the byte offset and byte size of the destination array to be a multiple of 4.","category":"page"},{"location":"faq/contributing/#Contributing","page":"Contributing","title":"Contributing","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Metal.jl is an especially accessible GPU backend with the presence of GPUs on Apple's recent popular Macbooks. As a result, an average Julia user can now develop and test GPU-accelerated code locally on their laptop. If you're using this package and see a bug or want some additional functionality, this page is for you. Hopefully this information helps encourage you to contribute to the package yourself.","category":"page"},{"location":"faq/contributing/#What-needs-help?","page":"Contributing","title":"What needs help?","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"If you didn't come to this page with your own feature to add, look at the current issues in the git repo for bugs and requested functionality.","category":"page"},{"location":"faq/contributing/#I'm-a-beginner,-can-I-help?","page":"Contributing","title":"I'm a beginner, can I help?","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Yes, but you may spend more time learning rather than directly contributing at the start. Depending on what your goals are though, this might be desirable. There are differing levels of difficulty when considering contributions to Metal.jl. If you're new to these things, check the issues for \"Good First Issue\" tags, look at the documentation for areas that could be added (beginners are especially good at detecting these sort of deficiencies), or message on the Slack #gpu channel asking for guidance.","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Regardless, if you've never used Metal.jl before, it'd probably be best to gain some exposure to it before trying to contibute. You might run into bugs yourself or discover some area you'd really like to help with.","category":"page"},{"location":"faq/contributing/#General-Workflow-for-Adding-Functionality","page":"Contributing","title":"General Workflow for Adding Functionality","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"If you're adding some functionality that originates from Metal Shading Language (MSL) directly (rather than high-level Julia functionality), the workflow will likely look like the below. If you're adding something that only relies on pure Julia additions, you will skip the first two steps.","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Create low-level, Julia wrappers for the Obj-C interface\nCreate high-level Julia structures and functionality\nCreate tests for added functionality","category":"page"},{"location":"faq/contributing/#Mapping-to-Metal-Intrinsics","page":"Contributing","title":"Mapping to Metal Intrinsics","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Some Metal functions map directly to Apple intermediate representation intrinsics. In this case, wrapping them into Metal.jl is relatively easy. All that needs to be done is to create a mapping from a Julia function via a simple ccall. See the threadgroup barrier implementation for reference.","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"However, the Metal documentation doesn't tell you what the format of the intrinsic names should be. To find this out, you need to create your own test kernel directly in the Metal Shading Language, compile it using Apple's tooling, then view the created intermediate representation (IR).","category":"page"},{"location":"faq/contributing/#Reverse-Engineering-Bare-MSL/Apple-IR","page":"Contributing","title":"Reverse-Engineering Bare MSL/Apple IR","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"First, you need to write an MSL kernel that uses the functionality you're interested in. For example,","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"#include \n\nusing namespace metal;\n\nkernel void dummy_kernel(device volatile atomic_float* out,\n uint i [[thread_position_in_grid]])\n{\n atomic_store_explicit(&out[i], 0.0f, memory_order_relaxed);\n}","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"To compile with Metal's tools and emit human-readable IR, run something roughly along the lines of: xcrun metal -S -emit-llvm dummy_kernel.metal","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"This will create a .ll file that you can then parse for whatever information you need. Be sure to double-check the metadata at the bottom for any significant changes your functionality introduces.","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Test with different types and configurations to see what changes are caused. Also ensure that when writing very simple kernels, whatever you're interested in doesn't get optimized away. Double-check that the kernel's IR makes sense for what you wrote.","category":"page"},{"location":"faq/contributing/#Metal-Performance-Shaders","page":"Contributing","title":"Metal Performance Shaders","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Metal exposes a special interface to its library of optimized kernels. Rather than accepting the normal set of input GPU data structures, it requires special MPS datatypes that assume row-major memory layout. As this is not the Julia default, adapt accordingly. Adding MPS functionality should be mostly straightforward, so this can be an easy entry point to helping. To get started, you can have a look at the Metal Performance Shaders Documentation from Apple.","category":"page"},{"location":"faq/contributing/#Exposing-your-Interface","page":"Contributing","title":"Exposing your Interface","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"There are varying degrees of user-facing interfaces from Metal.jl. At the lowest level is Metal.MTL.xxx. This is for low-level functionality close to or at bare Objective-C, or things that a normal user wouldn't directly be using. Metal.MPS.xxx is for Metal Performance Shader specifics (like MPSMatrix). Next, is Metal.xxx. This is for higher-level, usually pure-Julian functionality (like device()). The only thing beyond this is exporting into the global namespace. That would be useful for uniquely-named functions/structures/macros with clear and common use-cases (MtlArray or @metal).","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Additionally, you can override non-Metal.jl functions like LinearAlgebra.mul! seen here. This is essentially (ab)using multiple dispatch to specialize for certain cases (usually for more performant execution).","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"If your function is only available from within GPU kernels (like thread indexing intrinsics). Be sure to properly annotate with @device_function to ensure that calling from the host doesn't kill your Julia process.","category":"page"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Generally, think about how frequently you expect your addition to be used, how complex its use-case is, and whether or not it clashes/reimplements/optimizes existing functionality from outside Metal.jl. Put it behind the corresponding interface.","category":"page"},{"location":"faq/contributing/#Creating-Tests","page":"Contributing","title":"Creating Tests","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"As it's good practice, and JuliaGPU has great CI/CD workflows, your addition should have associated tests to ensure correctness and edge cases. Look to existing examples under the test folder for initial guidance, and be sure to create tests for all valid types. Any new Julia file in this folder will be ran as its own testset. If you feel your tests don't fit in any existing place, you'll probably want to create a new file with an appropriate name.","category":"page"},{"location":"faq/contributing/#Running-a-Subset-of-the-Existing-Tests","page":"Contributing","title":"Running a Subset of the Existing Tests","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Sometimes you won't want to run the entire testsuite. You may just want to run the tests for your new functionality. To do that, you can either pass the name of the testset to the test/runtests.jl script: julia --project=test test/runtests.jl metal or you can isolate test files by running them alone after running the test/setup.jl script: julia --project=test -L test/setup.jl test/metal.jl","category":"page"},{"location":"faq/contributing/#Thank-You-and-Good-Luck","page":"Contributing","title":"Thank You and Good Luck","text":"","category":"section"},{"location":"faq/contributing/","page":"Contributing","title":"Contributing","text":"Open-source projects like this only happen because people like you are willing to spend their free time helping out. Most anything you're able to do is helpful, but if you get stuck, seek guidance from Slack or Discourse. Don't feel like your contribution has to be perfect. If you put in effort and make progress, there will likely be some senior developer willing to polish your code before merging. Open-source software is a team effort...welcome to the team!","category":"page"},{"location":"api/kernel/#Kernel-programming","page":"Kernel programming","title":"Kernel programming","text":"","category":"section"},{"location":"api/kernel/","page":"Kernel programming","title":"Kernel programming","text":"This section lists the package's public functionality that corresponds to special Metal functions for use in device code. For more information about these functions, please consult the Metal Shading Language specification.","category":"page"},{"location":"api/kernel/","page":"Kernel programming","title":"Kernel programming","text":"This is made possible by interfacing with the Metal libraries by wrapping a subset of the ObjectiveC APIs using ObjectiveC.jl. These low-level wrappers are available in the MTL submodule exported by Metal.jl.","category":"page"},{"location":"api/kernel/#Indexing-and-dimensions","page":"Kernel programming","title":"Indexing and dimensions","text":"","category":"section"},{"location":"api/kernel/","page":"Kernel programming","title":"Kernel programming","text":"thread_execution_width\nthread_index_in_quadgroup\nthread_index_in_simdgroup\nthread_index_in_threadgroup\nthread_position_in_grid_1d\nthread_position_in_threadgroup_1d\nthreadgroup_position_in_grid_1d\nthreadgroups_per_grid_1d\nthreads_per_grid_1d\nthreads_per_simdgroup\nthreads_per_threadgroup_1d\nsimdgroups_per_threadgroup\nsimdgroup_index_in_threadgroup\nquadgroup_index_in_threadgroup\nquadgroups_per_threadgroup\ngrid_size_1d\ngrid_origin_1d","category":"page"},{"location":"api/kernel/#Metal.thread_execution_width","page":"Kernel programming","title":"Metal.thread_execution_width","text":"thread_execution_width()::UInt32\n\nReturn the execution width of the compute unit.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.thread_index_in_quadgroup","page":"Kernel programming","title":"Metal.thread_index_in_quadgroup","text":"thread_index_in_quadgroup()::UInt32\n\nReturn the index of the current thread in its quadgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.thread_index_in_simdgroup","page":"Kernel programming","title":"Metal.thread_index_in_simdgroup","text":"thread_index_in_simdgroup()::UInt32\n\nReturn the index of the current thread in its simdgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.thread_index_in_threadgroup","page":"Kernel programming","title":"Metal.thread_index_in_threadgroup","text":"thread_index_in_threadgroup()::UInt32\n\nReturn the index of the current thread in its threadgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.thread_position_in_grid_1d","page":"Kernel programming","title":"Metal.thread_position_in_grid_1d","text":"thread_position_in_grid_1d()::UInt32\nthread_position_in_grid_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}\nthread_position_in_grid_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}\n\nReturn the current thread's position in an N-dimensional grid of threads.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.thread_position_in_threadgroup_1d","page":"Kernel programming","title":"Metal.thread_position_in_threadgroup_1d","text":"thread_position_in_threadgroup_1d()::UInt32\nthread_position_in_threadgroup_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}\nthread_position_in_threadgroup_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}\n\nReturn the current thread's unique position within a threadgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.threadgroup_position_in_grid_1d","page":"Kernel programming","title":"Metal.threadgroup_position_in_grid_1d","text":"threadgroup_position_in_grid_1d()::UInt32\nthreadgroup_position_in_grid_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}\nthreadgroup_position_in_grid_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}\n\nReturn the current threadgroup's unique position within the grid.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.threadgroups_per_grid_1d","page":"Kernel programming","title":"Metal.threadgroups_per_grid_1d","text":"threadgroups_per_grid_1d()::UInt32\nthreadgroups_per_grid_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}\nthreadgroups_per_grid_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}\n\nReturn the number of threadgroups per grid.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.threads_per_grid_1d","page":"Kernel programming","title":"Metal.threads_per_grid_1d","text":"threads_per_grid_1d()::UInt32\nthreads_per_grid_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}\nthreads_per_grid_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}\n\nReturn the grid size.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.threads_per_simdgroup","page":"Kernel programming","title":"Metal.threads_per_simdgroup","text":"threads_per_simdgroup()::UInt32\n\nReturn the thread execution width of a simdgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.threads_per_threadgroup_1d","page":"Kernel programming","title":"Metal.threads_per_threadgroup_1d","text":"threads_per_threadgroup_1d()::UInt32\nthreads_per_threadgroup_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}\nthreads_per_threadgroup_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}\n\nReturn the thread execution width of a threadgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.simdgroups_per_threadgroup","page":"Kernel programming","title":"Metal.simdgroups_per_threadgroup","text":"simdgroups_per_threadgroup()::UInt32\n\nReturn the simdgroup execution width of a threadgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.simdgroup_index_in_threadgroup","page":"Kernel programming","title":"Metal.simdgroup_index_in_threadgroup","text":"simdgroup_index_in_threadgroup()::UInt32\n\nReturn the index of a simdgroup within a threadgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.quadgroup_index_in_threadgroup","page":"Kernel programming","title":"Metal.quadgroup_index_in_threadgroup","text":"quadgroup_index_in_threadgroup()::UInt32\n\nReturn the index of a quadgroup within a threadgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.quadgroups_per_threadgroup","page":"Kernel programming","title":"Metal.quadgroups_per_threadgroup","text":"quadgroups_per_threadgroup()::UInt32\n\nReturn the quadgroup execution width of a threadgroup.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.grid_size_1d","page":"Kernel programming","title":"Metal.grid_size_1d","text":"grid_size_1d()::UInt32\ngrid_size_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}\ngrid_size_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}\n\nReturn maximum size of the grid for threads that read per-thread stage-in data.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.grid_origin_1d","page":"Kernel programming","title":"Metal.grid_origin_1d","text":"grid_origin_1d()::UInt32\ngrid_origin_2d()::NamedTuple{(:x, :y), Tuple{UInt32, UInt32}}\ngrid_origin_3d()::NamedTuple{(:x, :y, :z), Tuple{UInt32, UInt32, UInt32}}\n\nReturn the origin offset of the grid for threads that read per-thread stage-in data.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Device-arrays","page":"Kernel programming","title":"Device arrays","text":"","category":"section"},{"location":"api/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Metal.jl provides a primitive, lightweight array type to manage GPU data organized in an plain, dense fashion. This is the device-counterpart to the MtlArray, and implements (part of) the array interface as well as other functionality for use on the GPU:","category":"page"},{"location":"api/kernel/","page":"Kernel programming","title":"Kernel programming","text":"MtlDeviceArray\nMetal.Const","category":"page"},{"location":"api/kernel/#Metal.MtlDeviceArray","page":"Kernel programming","title":"Metal.MtlDeviceArray","text":"MtlDeviceArray(dims, ptr)\nMtlDeviceArray{T}(dims, ptr)\nMtlDeviceArray{T,A}(dims, ptr)\nMtlDeviceArray{T,A,N}(dims, ptr)\n\nConstruct an N-dimensional dense Metal device array with element type T wrapping a pointer, where N is determined from the length of dims and T is determined from the type of ptr.\n\ndims may be a single scalar, or a tuple of integers corresponding to the lengths in each dimension). If the rank N is supplied explicitly as in Array{T,N}(dims), then it must match the length of dims. The same applies to the element type T, which should match the type of the pointer ptr.\n\n\n\n\n\n","category":"type"},{"location":"api/kernel/#Metal.Const","page":"Kernel programming","title":"Metal.Const","text":"Const(A::MtlDeviceArray)\n\nMark a MtlDeviceArray as constant/read-only and to use the constant address space.\n\nwarning: Warning\nExperimental API. Subject to change without deprecation.\n\n\n\n\n\n","category":"type"},{"location":"api/kernel/#Shared-memory","page":"Kernel programming","title":"Shared memory","text":"","category":"section"},{"location":"api/kernel/","page":"Kernel programming","title":"Kernel programming","text":"MtlThreadGroupArray","category":"page"},{"location":"api/kernel/#Metal.MtlThreadGroupArray","page":"Kernel programming","title":"Metal.MtlThreadGroupArray","text":"MtlThreadGroupArray(::Type{T}, dims)\n\nCreate an array local to each threadgroup launched during kernel execution.\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Synchronization","page":"Kernel programming","title":"Synchronization","text":"","category":"section"},{"location":"api/kernel/","page":"Kernel programming","title":"Kernel programming","text":"MemoryFlags\nthreadgroup_barrier\nsimdgroup_barrier","category":"page"},{"location":"api/kernel/#Metal.MemoryFlags","page":"Kernel programming","title":"Metal.MemoryFlags","text":"MemoryFlags\n\nFlags to set the memory synchronization behavior of threadgroup_barrier and simdgroup_barrier.\n\nPossible values:\n\nNone: Set barriers to only act as an execution barrier and not apply a memory fence.\n\nDevice: Ensure the GPU correctly orders the memory operations to device memory\n for threads in the threadgroup or simdgroup.\n\nThreadGroup: Ensure the GPU correctly orders the memory operations to threadgroup\n memory for threads in a threadgroup or simdgroup.\n\nTexture: Ensure the GPU correctly orders the memory operations to texture memory for\n threads in a threadgroup or simdgroup for a texture with the read_write access qualifier.\n\nThreadGroup_ImgBlock: Ensure the GPU correctly orders the memory operations to threadgroup imageblock memory\n for threads in a threadgroup or simdgroup.\n\n\n\n\n\n","category":"type"},{"location":"api/kernel/#Metal.threadgroup_barrier","page":"Kernel programming","title":"Metal.threadgroup_barrier","text":"threadgroup_barrier(flag::MemoryFlags=MemoryFlagNone)\n\nSynchronize all threads in a threadgroup.\n\nPossible flags that affect the memory synchronization behavior are found in MemoryFlags\n\n\n\n\n\n","category":"function"},{"location":"api/kernel/#Metal.simdgroup_barrier","page":"Kernel programming","title":"Metal.simdgroup_barrier","text":"simdgroup_barrier(flag::MemoryFlags=MemoryFlagNone)\n\nSynchronize all threads in a SIMD-group.\n\nPossible flags that affect the memory synchronization behavior are found in MemoryFlags\n\n\n\n\n\n","category":"function"},{"location":"faq/faq/#Frequently-Asked-Questions","page":"Frequently Asked Questions","title":"Frequently Asked Questions","text":"","category":"section"},{"location":"faq/faq/#Can-you-wrap-this-Metal-API?","page":"Frequently Asked Questions","title":"Can you wrap this Metal API?","text":"","category":"section"},{"location":"faq/faq/","page":"Frequently Asked Questions","title":"Frequently Asked Questions","text":"Most likely. Any help on designing or implementing high-level wrappers for MSL's low-level functionality is greatly appreciated, so please consider contributing your uses of these APIs on the respective repositories.","category":"page"},{"location":"api/mps/#Metal-Performance-Shaders","page":"Metal Performance Shaders","title":"Metal Performance Shaders","text":"","category":"section"},{"location":"api/mps/","page":"Metal Performance Shaders","title":"Metal Performance Shaders","text":"This section lists the package's public functionality that corresponds to the Metal Performance Shaders functions. For more information about these functions, or to see which functions have yet to be implemented in this package, please consult the Metal Performance Shaders Documentation.","category":"page"},{"location":"api/mps/#Matrices-and-Vectors","page":"Metal Performance Shaders","title":"Matrices and Vectors","text":"","category":"section"},{"location":"api/mps/","page":"Metal Performance Shaders","title":"Metal Performance Shaders","text":"MPS.MPSMatrix\nMPS.MPSVector","category":"page"},{"location":"api/mps/#Metal.MPS.MPSMatrix","page":"Metal Performance Shaders","title":"Metal.MPS.MPSMatrix","text":"MPSMatrix(mat::MtlMatrix)\n\nMetal matrix representation used in Performance Shaders.\n\nNote that this results in a transposed view of the input, as Metal stores matrices row-major instead of column-major.\n\n\n\n\n\nMPSMatrix(vec::MtlVector)\n\nMetal matrix representation used in Performance Shaders.\n\nNote that this results in a transposed view of the input, as Metal stores matrices row-major instead of column-major.\n\n\n\n\n\nMPSMatrix(arr::MtlArray{T,3})\n\nMetal batched matrix representation used in Performance Shaders.\n\nNote that this results in a transposed view of the input, as Metal stores matrices row-major instead of column-major.\n\n\n\n\n\n","category":"type"},{"location":"api/mps/#Metal.MPS.MPSVector","page":"Metal Performance Shaders","title":"Metal.MPS.MPSVector","text":"MPSVector(arr::MtlVector)\n\nMetal vector representation used in Performance Shaders.\n\n\n\n\n\n","category":"type"},{"location":"api/mps/#Matrix-Arithmetic-Operators","page":"Metal Performance Shaders","title":"Matrix Arithmetic Operators","text":"","category":"section"},{"location":"api/mps/","page":"Metal Performance Shaders","title":"Metal Performance Shaders","text":"MPS.matmul!\nMPS.matvecmul!\nMPS.topk\nMPS.topk!","category":"page"},{"location":"api/mps/#Metal.MPS.matmul!","page":"Metal Performance Shaders","title":"Metal.MPS.matmul!","text":"matMulMPS(a::MtlMatrix, b::MtlMatrix, c::MtlMatrix, alpha=1, beta=1,\n transpose_left=false, transpose_right=false)\n\nA MPSMatrixMultiplication kernel thay computes: c = alpha * op(a) * beta * op(b) + beta * C\n\nThis function should not typically be used. Rather, use the normal LinearAlgebra interface with any MtlArray and it should be accelerated using Metal Performance Shaders.\n\n\n\n\n\n","category":"function"},{"location":"api/mps/#Metal.MPS.matvecmul!","page":"Metal Performance Shaders","title":"Metal.MPS.matvecmul!","text":"matvecmul!(c::MtlVector, a::MtlMatrix, b::MtlVector, alpha=1, beta=1, transpose=false)\n\nA MPSMatrixVectorMultiplication kernel thay computes: c = alpha * op(a) * b + beta * c\n\nThis function should not typically be used. Rather, use the normal LinearAlgebra interface with any MtlArray and it should be accelerated using Metal Performance Shaders.\n\n\n\n\n\n","category":"function"},{"location":"api/mps/#Metal.MPS.topk","page":"Metal Performance Shaders","title":"Metal.MPS.topk","text":"MPS.topk(A::MtlMatrix{T}, k) where {T<:MtlFloat}\n\nCompute the top k values and their corresponding indices column-wise in a matrix A. Return the indices in I and the values in V.\n\nk cannot be greater than 16.\n\nUses MPSMatrixFindTopK.\n\nSee also: topk!.\n\nwarn: Warn\nThis interface is experimental, and might change without warning.\n\n\n\n\n\n","category":"function"},{"location":"api/mps/#Metal.MPS.topk!","page":"Metal Performance Shaders","title":"Metal.MPS.topk!","text":"MPS.topk!(A::MtlMatrix{T}, I::MtlMatrix{Int32}, V::MtlMatrix{T}, k)\n where {T<:MtlFloat}\n\nCompute the top k values and their corresponding indices column-wise in a matrix A. Return the indices in I and the values in V.\n\nk cannot be greater than 16.\n\nUses MPSMatrixFindTopK.\n\nSee also: topk.\n\nwarn: Warn\nThis interface is experimental, and might change without warning.\n\n\n\n\n\n","category":"function"},{"location":"api/array/#Array-programming","page":"Array programming","title":"Array programming","text":"","category":"section"},{"location":"api/array/","page":"Array programming","title":"Array programming","text":"The Metal array type, MtlArray, generally implements the Base array interface and all of its expected methods.","category":"page"},{"location":"api/array/","page":"Array programming","title":"Array programming","text":"However, there is the special function mtl for transferring an array over to the gpu. For compatibility reasons, it will automatically convert arrays of Float64 to Float32.","category":"page"},{"location":"api/array/","page":"Array programming","title":"Array programming","text":"mtl\nMtlArray\nMtlVector\nMtlMatrix\nMtlVecOrMat","category":"page"},{"location":"api/array/#Metal.mtl","page":"Array programming","title":"Metal.mtl","text":"mtl(A; storage=Metal.PrivateStorage)\n\nstorage can be Metal.PrivateStorage (default), Metal.SharedStorage, or Metal.ManagedStorage.\n\nOpinionated GPU array adaptor, which may alter the element type T of arrays:\n\nFor T<:AbstractFloat, it makes a MtlArray{Float32} for performance and compatibility reasons (except for Float16).\nFor T<:Complex{<:AbstractFloat} it makes a MtlArray{ComplexF32}.\nFor other isbitstype(T), it makes a MtlArray{T}.\n\nBy contrast, MtlArray(A) never changes the element type.\n\nUses Adapt.jl to act inside some wrapper structs.\n\nExamples\n\njulia> mtl(ones(3)')\n1×3 adjoint(::MtlVector{Float32, Metal.PrivateStorage}) with eltype Float32:\n 1.0 1.0 1.0\n\njulia> mtl(zeros(1,3); storage=Metal.SharedStorage)\n1×3 MtlMatrix{Float32, Metal.SharedStorage}:\n 0.0 0.0 0.0\n\njulia> mtl(1:3)\n1:3\n\njulia> MtlArray(1:3)\n3-element MtlVector{Int64, Metal.PrivateStorage}:\n 1\n 2\n 3\n\n\n\n\n\n","category":"function"},{"location":"api/array/#Metal.MtlArray","page":"Array programming","title":"Metal.MtlArray","text":"MtlArray{T,N,S} <: AbstractGPUArray{T,N}\n\nN-dimensional Metal array with storage mode S and elements of type T.\n\nS can be Metal.PrivateStorage (default), Metal.SharedStorage, or Metal.ManagedStorage.\n\nSee the Array Programming section of the Metal.jl docs for more details.\n\n\n\n\n\n","category":"type"},{"location":"api/array/#Metal.MtlVector","page":"Array programming","title":"Metal.MtlVector","text":"MtlVector{T,S} <: AbstractGPUVector{T}\n\nOne-dimensional array with elements of type T for use with Apple Metal-compatible GPUs. Alias for MtlArray{T,1,S}.\n\nSee also Vector(@ref), and the Array Programming section of the Metal.jl docs for more details.\n\n\n\n\n\n","category":"type"},{"location":"api/array/#Metal.MtlMatrix","page":"Array programming","title":"Metal.MtlMatrix","text":"MtlMatrix{T,S} <: AbstractGPUMatrix{T}\n\nTwo-dimensional array with elements of type T for use with Apple Metal-compatible GPUs. Alias for MtlArray{T,2,S}.\n\nSee also Matrix(@ref), and the Array Programming section of the Metal.jl docs for more details.\n\n\n\n\n\n","category":"type"},{"location":"api/array/#Metal.MtlVecOrMat","page":"Array programming","title":"Metal.MtlVecOrMat","text":"MtlVecOrMat{T,S}\n\nUnion type of MtlVector{T,S} and MtlMatrix{T,S} which allows functions to accept either an MtlMatrix or an MtlVector.\n\nSee also VecOrMat(@ref) for examples.\n\n\n\n\n\n","category":"type"},{"location":"api/array/#Storage-modes","page":"Array programming","title":"Storage modes","text":"","category":"section"},{"location":"api/array/","page":"Array programming","title":"Array programming","text":"The Metal API has various storage modes that dictate how a resource can be accessed. MtlArrays are Metal.PrivateStorage by default, but they can also be Metal.SharedStorage or Metal.ManagedStorage. For more information on storage modes, see the official Metal documentation.","category":"page"},{"location":"api/array/","page":"Array programming","title":"Array programming","text":"Metal.PrivateStorage\nMetal.SharedStorage\nMetal.ManagedStorage","category":"page"},{"location":"api/array/#Metal.MTL.PrivateStorage","page":"Array programming","title":"Metal.MTL.PrivateStorage","text":"struct Metal.PrivateStorage <: MTL.StorageMode\n\nUsed to indicate that the resource is stored using MTLStorageModePrivate in memory.\n\nFor more information on Metal storage modes, refer to the official Metal documentation.\n\nSee also Metal.SharedStorage and Metal.ManagedStorage.\n\n\n\n\n\n","category":"type"},{"location":"api/array/#Metal.MTL.SharedStorage","page":"Array programming","title":"Metal.MTL.SharedStorage","text":"struct Metal.SharedStorage <: MTL.StorageMode\n\nUsed to indicate that the resource is stored using MTLStorageModeShared in memory.\n\nFor more information on Metal storage modes, refer to the official Metal documentation.\n\nSee also Metal.PrivateStorage and Metal.ManagedStorage.\n\n\n\n\n\n","category":"type"},{"location":"api/array/#Metal.MTL.ManagedStorage","page":"Array programming","title":"Metal.MTL.ManagedStorage","text":"struct Metal.ManagedStorage <: MTL.StorageMode\n\nUsed to indicate that the resource is stored using MTLStorageModeManaged in memory.\n\nFor more information on Metal storage modes, refer to the official Metal documentation.\n\nSee also Metal.SharedStorage and Metal.PrivateStorage.\n\n\n\n\n\n","category":"type"},{"location":"api/array/","page":"Array programming","title":"Array programming","text":"There also exist the following convenience functions to check if an MtlArray is using a specific storage mode:","category":"page"},{"location":"api/array/","page":"Array programming","title":"Array programming","text":"is_private\nis_shared\nis_managed","category":"page"},{"location":"api/array/#Metal.is_private","page":"Array programming","title":"Metal.is_private","text":"is_private(A::MtlArray) -> Bool\n\nReturns true if A has storage mode Metal.PrivateStorage.\n\nSee also is_shared and is_managed.\n\n\n\n\n\n","category":"function"},{"location":"api/array/#Metal.is_shared","page":"Array programming","title":"Metal.is_shared","text":"is_shared(A::MtlArray) -> Bool\n\nReturns true if A has storage mode Metal.SharedStorage.\n\nSee also is_private and is_managed.\n\n\n\n\n\n","category":"function"},{"location":"api/array/#Metal.is_managed","page":"Array programming","title":"Metal.is_managed","text":"is_managed(A::MtlArray) -> Bool\n\nReturns true if A has storage mode Metal.ManagedStorage.\n\nSee also is_shared and is_private.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Essentials","page":"Essentials","title":"Essentials","text":"","category":"section"},{"location":"api/essentials/#Versions-and-Support","page":"Essentials","title":"Versions and Support","text":"","category":"section"},{"location":"api/essentials/","page":"Essentials","title":"Essentials","text":"Metal.macos_version\nMetal.darwin_version\nMetal.metal_support\nMetal.metallib_support\nMetal.air_support","category":"page"},{"location":"api/essentials/#Metal.macos_version","page":"Essentials","title":"Metal.macos_version","text":"Metal.macos_version() -> VersionNumber\n\nReturns the host macOS version.\n\nSee also Metal.darwin_version.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.darwin_version","page":"Essentials","title":"Metal.darwin_version","text":"Metal.darwin_version() -> VersionNumber\n\nReturns the host Darwin kernel version.\n\nSee also Metal.macos_version.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.metal_support","page":"Essentials","title":"Metal.metal_support","text":"Metal.metal_support() -> VersionNumber\n\nReturns the highest supported version for the Metal Shading Language.\n\nSee also Metal.metallib_support and Metal.air_support.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.metallib_support","page":"Essentials","title":"Metal.metallib_support","text":"Metal.metallib_support() -> VersionNumber\n\nReturns the highest supported version for the metallib file format.\n\nSee also Metal.air_support and Metal.metal_support.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.air_support","page":"Essentials","title":"Metal.air_support","text":"Metal.air_support() -> VersionNumber\n\nReturns the highest supported version for the embedded AIR bitcode format.\n\nSee also Metal.metallib_support and Metal.metal_support.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Global-State","page":"Essentials","title":"Global State","text":"","category":"section"},{"location":"api/essentials/","page":"Essentials","title":"Essentials","text":"Metal.device!\nMetal.devices\nMetal.device\nMetal.global_queue\nMetal.synchronize\nMetal.device_synchronize","category":"page"},{"location":"api/essentials/#Metal.device!","page":"Essentials","title":"Metal.device!","text":"device!(dev::MTLDevice)\n\nSets the Metal GPU device associated with the current Julia task.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.MTL.devices","page":"Essentials","title":"Metal.MTL.devices","text":"devices()\n\nGet an iterator for the compute devices.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.device","page":"Essentials","title":"Metal.device","text":"device()::MTLDevice\n\nReturn the Metal GPU device associated with the current Julia task.\n\nSince all M-series systems currently only externally show a single GPU, this function effectively returns the only system GPU.\n\n\n\n\n\ndevice(<:MtlArray)\n\nGet the Metal device for an MtlArray.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.global_queue","page":"Essentials","title":"Metal.global_queue","text":"global_queue(dev::MTLDevice)::MTLCommandQueue\n\nReturn the Metal command queue associated with the current Julia thread.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.synchronize","page":"Essentials","title":"Metal.synchronize","text":"synchronize(queue)\n\nWait for currently committed GPU work on this queue to finish.\n\nCreate a new MTLCommandBuffer from the global command queue, commit it to the queue, and simply wait for it to be completed. Since command buffers should execute in a First-In-First-Out manner, this synchronizes the GPU.\n\n\n\n\n\n","category":"function"},{"location":"api/essentials/#Metal.device_synchronize","page":"Essentials","title":"Metal.device_synchronize","text":"device_synchronize()\n\nSynchronize all committed GPU work across all global queues\n\n\n\n\n\n","category":"function"},{"location":"usage/kernel/#Kernel-programming","page":"Kernel programming","title":"Kernel programming","text":"","category":"section"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Metal.jl is based off of Apple's Metal Shading Language (MSL) and Metal framework. The interface allows you to utilize the graphics and computing power of Mac GPUs. Like many other GPU frameworks, its history is rooted in graphics processing but has found use in computing/general purpose GPU (GPGPU) applications.","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"The most fundamental idea of programming GPUs (when compared to serial CPU programming) is its parallelism. A GPU function (kernel), when called, is not just ran once in isolation. Rather, numerous (often thousands to millions) psuedo-independent instances (called threads) of the kernel are executed in parallel. These threads are arranged in a hierarchy that allows for varying levels of synchronization. For Metal, the hierarchy is as follows:","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Thread: A single execution unit of the kernel\nThreadgroup: A collection of threads that share a common block of memory and synchronization","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"barriers","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Grid: A collection of threadgroups","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"The threadgroup and grid sizes are set by the user when launching the GPU kernel. There are upper limits determined by the targeted hardware, and the sizes can be 1, 2, or 3-dimensional. For Metal.jl, these sizes are set using the @metal macro's keyword arguments. The grid keyword determines the grid size while the threads keyword determines the threadgroup size.","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"For example, given a 10x10x3 image that you want to run a function independently on each pixel, the kernel launch code might look like the following: @metal threads=(10,10) groups=3 my_kernel(gpu_image_array) This would launch 3 separate threadgroups of 100 threads each (10 in the first dimension and 10 in the second dimension)","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"There is also additional hierarchy layers that consists of small groups of threads that execute in lockstep called waves/SIMD groups/wavefronts* and quadgroups. However, the basic three-tier hierarchy is enough to get started.","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Here is a helpful link with good visualizations of Metal's thread hierarchy (also covering SIMD groups).","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Each thread has its own set of private variables. Most importantly, each thread has associated unique indices to identify itself within its threadgroup and grid. These are traditionally what are used to differentiate execution across threads. You can also query what the grid and threadgroup sizes are as well.","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"For Metal.jl, these values are accessed via the following functions:","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"thread_index_in_threadgroup()\ngrid_size_Xd()\nthread_position_in_grid_Xd()\nthread_position_in_threadgroup_Xd()\nthreadgroup_position_in_grid_Xd()\nthreadgroups_per_grid_Xd()\nthreads_per_grid_Xd()\nthreads_per_threadgroup_Xd()","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Where 'X' is 1, 2, or 3 according to the number of dimensions requested.","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Using these in a kernel (taken directly from the vadd example):","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"function vadd(a, b, c)\n i = thread_position_in_grid_1d()\n c[i] = a[i] + b[i]\n return\nend","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"This kernel takes in three vectors (a,b,c) all of the same length and stores the element-wise sum of a and b into c. Each thread in this kernel gets its unique position in the grid (arrangement of all threadgroups) and stores this value into the variable i which is then used as the index into the vectors. Thus, each thread is computing one sum and storing the result in the output vector.","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"To ensure this kernel functions properly, we have to launch it with exactly as many threads as the length of the vectors. If we under or over-launch threads, the result could be incorrect.","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"An example of a good launch:","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"len = prod(size(d_a))\n@metal threads=len vadd(d_a, d_b, d_c)","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Additional notes:","category":"page"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Kernels must always return nothing\nKernels are asynchronous. To synchronize, use the Metal.@sync macro.","category":"page"},{"location":"usage/kernel/#Other-Helpful-Links","page":"Kernel programming","title":"Other Helpful Links","text":"","category":"section"},{"location":"usage/kernel/","page":"Kernel programming","title":"Kernel programming","text":"Metal Shading Language Specification","category":"page"},{"location":"api/compiler/#Compiler","page":"Compiler","title":"Compiler","text":"","category":"section"},{"location":"api/compiler/#Execution","page":"Compiler","title":"Execution","text":"","category":"section"},{"location":"api/compiler/","page":"Compiler","title":"Compiler","text":"The main entry-point to the compiler is the @metal macro:","category":"page"},{"location":"api/compiler/","page":"Compiler","title":"Compiler","text":"@metal","category":"page"},{"location":"api/compiler/#Metal.@metal","page":"Compiler","title":"Metal.@metal","text":"@metal threads=... groups=... [kwargs...] func(args...)\n\nHigh-level interface for executing code on a GPU.\n\nThe @metal macro should prefix a call, with func a callable function or object that should return nothing. It will be compiled to a Metal function upon first use, and to a certain extent arguments will be converted and managed automatically using mtlconvert. Finally, a call to mtlcall is performed, creating a command buffer in the current global command queue then committing it.\n\nThere is one supported keyword argument that influences the behavior of @metal:\n\nlaunch: whether to launch this kernel, defaults to true. If false the returned kernel object should be launched by calling it and passing arguments again.\nname: the name of the kernel in the generated code. Defaults to an automatically- generated name.\nqueue: the command queue to use for this kernel. Defaults to the global command queue.\n\n\n\n\n\n","category":"macro"},{"location":"api/compiler/","page":"Compiler","title":"Compiler","text":"If needed, you can use a lower-level API that lets you inspect the compiler kernel:","category":"page"},{"location":"api/compiler/","page":"Compiler","title":"Compiler","text":"Metal.mtlconvert\nMetal.mtlfunction","category":"page"},{"location":"api/compiler/#Metal.mtlconvert","page":"Compiler","title":"Metal.mtlconvert","text":"mtlconvert(x, [cce])\n\nThis function is called for every argument to be passed to a kernel, allowing it to be converted to a GPU-friendly format. By default, the function does nothing and returns the input object x as-is.\n\nDo not add methods to this function, but instead extend the underlying Adapt.jl package and register methods for the the Metal.Adaptor type.\n\n\n\n\n\n","category":"function"},{"location":"api/compiler/#Metal.mtlfunction","page":"Compiler","title":"Metal.mtlfunction","text":"mtlfunction(f, tt=Tuple{}; kwargs...)\n\nLow-level interface to compile a function invocation for the currently-active GPU, returning a callable kernel object. For a higher-level interface, use @metal.\n\nThe output of this function is automatically cached, i.e. you can simply call mtlfunction in a hot path without degrading performance. New code will be generated automatically when the function changes, or when different types or keyword arguments are provided.\n\n\n\n\n\n","category":"function"},{"location":"api/compiler/#Reflection","page":"Compiler","title":"Reflection","text":"","category":"section"},{"location":"api/compiler/","page":"Compiler","title":"Compiler","text":"If you want to inspect generated code, you can use macros that resemble functionality from the InteractiveUtils standard library:","category":"page"},{"location":"api/compiler/","page":"Compiler","title":"Compiler","text":"@device_code_lowered\n@device_code_typed\n@device_code_warntype\n@device_code_llvm\n@device_code_native\n@device_code_agx\n@device_code","category":"page"},{"location":"api/compiler/","page":"Compiler","title":"Compiler","text":"For more information, please consult the GPUCompiler.jl documentation. code_agx is actually code_native:","category":"page"},{"location":"#MacOS-GPU-programming-in-Julia","page":"Home","title":"MacOS GPU programming in Julia","text":"","category":"section"},{"location":"","page":"Home","title":"Home","text":"The Metal.jl package is the main entry point for GPU programming on MacOS in Julia. The package makes it possible to do so at various abstraction levels, from easy-to-use arrays down to hand-written kernels using low-level Metal APIs.","category":"page"},{"location":"","page":"Home","title":"Home","text":"If you have any questions, please feel free to use the #gpu channel on the Julia slack, or the GPU domain of the Julia Discourse.","category":"page"},{"location":"","page":"Home","title":"Home","text":"As this package is still under development, if you spot a bug, please file an issue.","category":"page"},{"location":"#Quick-Start","page":"Home","title":"Quick Start","text":"","category":"section"},{"location":"","page":"Home","title":"Home","text":"Metal.jl ties into your system's existing Metal Shading Language compiler toolchain, so no additional installs are required (unless you want to view profiled GPU operations)","category":"page"},{"location":"","page":"Home","title":"Home","text":"# install the package\nusing Pkg\nPkg.add(\"Metal\")\n\n# smoke test\nusing Metal\nMetal.versioninfo()","category":"page"},{"location":"","page":"Home","title":"Home","text":"If you want to ensure everything works as expected, you can execute the test suite.","category":"page"},{"location":"","page":"Home","title":"Home","text":"using Pkg\nPkg.test(\"Metal\")","category":"page"},{"location":"","page":"Home","title":"Home","text":"The following resources may also be of interest (although are mainly focused on the CUDA GPU backend):","category":"page"},{"location":"","page":"Home","title":"Home","text":"Effectively using GPUs with Julia: slides\nHow Julia is compiled to GPUs: video","category":"page"},{"location":"#Contributing","page":"Home","title":"Contributing","text":"","category":"section"},{"location":"","page":"Home","title":"Home","text":"If you want to help improve this package, look at the contributing page for more details.","category":"page"},{"location":"#Acknowledgements","page":"Home","title":"Acknowledgements","text":"","category":"section"},{"location":"","page":"Home","title":"Home","text":"The Julia Metal stack has been a collaborative effort by many individuals. Significant contributions have been made by the following individuals:","category":"page"},{"location":"","page":"Home","title":"Home","text":"Tim Besard (@maleadt) (lead developer)\nFilippo Vicentini (@PhilipVinc)\nMax Hawkins (@max-Hawkins)","category":"page"},{"location":"#Supporting-and-Citing","page":"Home","title":"Supporting and Citing","text":"","category":"section"},{"location":"","page":"Home","title":"Home","text":"Some of the software in this ecosystem was developed as part of academic research. If you would like to help support it, please star the repository as such metrics may help us secure funding in the future. If you use our software as part of your research, teaching, or other activities, we would be grateful if you could cite our work. The CITATION.cff file in the root of this repository lists the relevant papers.","category":"page"},{"location":"profiling/#Profiling","page":"Profiling","title":"Profiling","text":"","category":"section"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"Profiling GPU code is harder than profiling Julia code executing on the CPU. For one, kernels typically execute asynchronously, and thus require appropriate synchronization when measuring their execution time. Furthermore, because the code executes on a different processor, it is much harder to know what is currently executing.","category":"page"},{"location":"profiling/#Time-measurements","page":"Profiling","title":"Time measurements","text":"","category":"section"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"For robust measurements, it is advised to use the BenchmarkTools.jl package which goes to great lengths to perform accurate measurements. Due to the asynchronous nature of GPUs, you need to ensure the GPU is synchronized at the end of every sample, e.g. by calling synchronize() or, even better, wrapping your code in Metal.@sync:","category":"page"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"Note that the allocations as reported by BenchmarkTools are CPU allocations.","category":"page"},{"location":"profiling/#Application-tracing","page":"Profiling","title":"Application tracing","text":"","category":"section"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"For profiling large applications, simple timings are insufficient. Instead, we want an overview of how and when the GPU was active to avoid times where the device was idle and/or find which kernels needs optimization.","category":"page"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"As we cannot use the Julia profiler for this task, we will use Metal's GPU profiler directly. Use the Metal.@profile macro to surround the code code of interest. This macro tells your system to track GPU calls and usage statistics and will save this information in a temporary folder ending in '.trace'. For later viewing, copy this folder to a stable location or use the 'dir' argument of the profile macro to store the gputrace to a different location directly.","category":"page"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"The resulting trace can be opened with the Instruments app, part of Xcode.","category":"page"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"julia> using Metal\n\njulia> function vadd(a, b, c)\n i = thread_position_in_grid_1d()\n c[i] = a[i] + b[i]\n return\n end\njulia> a = MtlArray([1]); b = MtlArray([2]); c = similar(a);\n\njulia> Metal.@profile @metal threads=length(c) vadd(a, b, c);\n...\n[ Info: System trace saved to julia_3.trace; open the resulting trace in Instruments","category":"page"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"It is possible to augment the trace with additional information by using signposts: Similar to NVTX markers and ranges in CUDA.jl, signpost intervals and events can be used to add respectively time intervals and points of interest to the trace. This can be done by using the signpost functionality from ObjectiveC.jl:","category":"page"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"using ObjectiveC, .OS\n\n@signpost_interval \"My Interval\" begin\n # code to profile\n @signpost_event \"My Event\"\nend","category":"page"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"For more information, e.g. how to pass additional messages to the signposts, or how to use a custom logger, consult the ObjectiveC.jl documentation, or the docstrings of the @signpost_interval and @signpost_event macros.","category":"page"},{"location":"profiling/#Frame-capture","page":"Profiling","title":"Frame capture","text":"","category":"section"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"For more details on specific operations, you can use Metal's frame capture feature to generate a more detailed, and replayable trace of the GPU operations. This requires that Julia is started with the METAL_CAPTURE_ENABLED environment variable set to 1. Frames are captured by wrapping the code of interest in Metal.@capture, and the resulting trace can be opened with Xcode.","category":"page"},{"location":"profiling/","page":"Profiling","title":"Profiling","text":"$ METAL_CAPTURE_ENABLED=1 julia\n...\n\njulia> using Metal\n\njulia> function vadd(a, b, c)\n i = thread_position_in_grid_1d()\n c[i] = a[i] + b[i]\n return\n end\n\njulia> a = MtlArray([1]); b = MtlArray([2]); c = similar(a);\n... Metal GPU Frame Capture Enabled\n\njulia> Metal.@capture @metal threads=length(c) vadd(a, b, c);\n...\n[ Info: GPU frame capture saved to julia_1.gputrace; open the resulting trace in Xcode","category":"page"}] } diff --git a/dev/usage/array/index.html b/dev/usage/array/index.html index e9dbe439..b7d0f531 100644 --- a/dev/usage/array/index.html +++ b/dev/usage/array/index.html @@ -1,5 +1,5 @@ -Array programming · Metal.jl

Array programming

The easiest way to use the GPU's massive parallelism, is by expressing operations in terms of arrays: Metal.jl provides an array type, MtlArray, and many specialized array operations that execute efficiently on the GPU hardware. In this section, we will briefly demonstrate use of the MtlArray type. Since we expose Metal's functionality by implementing existing Julia interfaces on the MtlArray type, you should refer to the upstream Julia documentation for more information on these operations.

If you encounter missing functionality, or are running into operations that trigger so-called "scalar iteration", have a look at the issue tracker and file a new issue if there's none. Do note that you can always access the underlying Metal APIs by calling into the relevant submodule.

Construction and Initialization

The MtlArray type aims to implement the AbstractArray interface, and provide implementations of methods that are commonly used when working with arrays. That means you can construct MtlArrays in the same way as regular Array objects:

julia> MtlArray{Int}(undef, 2)
+Array programming · Metal.jl
Array programming
The easiest way to use the GPU's massive parallelism, is by expressing operations in terms of arrays: Metal.jl provides an array type, MtlArray, and many specialized array operations that execute efficiently on the GPU hardware. In this section, we will briefly demonstrate use of the MtlArray type. Since we expose Metal's functionality by implementing existing Julia interfaces on the MtlArray type, you should refer to the upstream Julia documentation for more information on these operations.
If you encounter missing functionality, or are running into operations that trigger so-called "scalar iteration", have a look at the issue tracker and file a new issue if there's none. Do note that you can always access the underlying Metal APIs by calling into the relevant submodule.
Construction and Initialization
The MtlArray type aims to implement the AbstractArray interface, and provide implementations of methods that are commonly used when working with arrays. That means you can construct MtlArrays in the same way as regular Array objects:
julia> MtlArray{Int}(undef, 2)
 2-element MtlVector{Int64, Metal.PrivateStorage}:
  0
  0
@@ -50,4 +50,20 @@
 
 julia> Base.mapreducedim!(identity, +, b, a)
 1×1 MtlMatrix{Float32, Metal.PrivateStorage}:
- 6.0

+ 6.0

Random numbers

Base's convenience functions for generating random numbers are available in Metal as well:

julia> Metal.rand(2)
+2-element MtlVector{Float32, Metal.PrivateStorage}:
+ 0.89025915
+ 0.8946847
+
+julia> Metal.randn(Float32, 2, 1)
+2×1 MtlMatrix{Float32, Metal.PrivateStorage}:
+ 1.2279074
+ 1.2518331

Behind the scenes, these random numbers come from two different generators: one backed by Metal Performance Shaders, another by using the GPUArrays.jl random methods. Operations on these generators are implemented using methods from the Random standard library:

julia> using Random, GPUArrays
+
+julia> a = Random.rand(MPS.default_rng(), Float32, 1)
+1-element MtlVector{Float32, Metal.PrivateStorage}:
+ 0.89025915
+
+julia> a = Random.rand!(GPUArrays.default_rng(MtlArray), a)
+1-element MtlVector{Float32, Metal.PrivateStorage}:
+ 0.0705002
Note

MPSMatrixRandom functionality requires Metal.jl >= v1.4

Warning

Random.rand!(::MPS.RNG, args...) and Random.randn!(::MPS.RNG, args...) have a framework limitation that requires the byte offset and byte size of the destination array to be a multiple of 4.

diff --git a/dev/usage/kernel/index.html b/dev/usage/kernel/index.html index a5e41a1a..e1ffdf66 100644 --- a/dev/usage/kernel/index.html +++ b/dev/usage/kernel/index.html @@ -4,4 +4,4 @@ c[i] = a[i] + b[i] return end

This kernel takes in three vectors (a,b,c) all of the same length and stores the element-wise sum of a and b into c. Each thread in this kernel gets its unique position in the grid (arrangement of all threadgroups) and stores this value into the variable i which is then used as the index into the vectors. Thus, each thread is computing one sum and storing the result in the output vector.

To ensure this kernel functions properly, we have to launch it with exactly as many threads as the length of the vectors. If we under or over-launch threads, the result could be incorrect.

An example of a good launch:

len = prod(size(d_a))
-@metal threads=len vadd(d_a, d_b, d_c)

Additional notes:

Metal Shading Language Specification An Introduction to GPU Programming course from University of Illinois (primarily in CUDA, but the concepts are transferable)

+@metal threads=len vadd(d_a, d_b, d_c)

Additional notes:

Metal Shading Language Specification

diff --git a/dev/usage/overview/index.html b/dev/usage/overview/index.html index bba0cb8a..514f973f 100644 --- a/dev/usage/overview/index.html +++ b/dev/usage/overview/index.html @@ -9,4 +9,4 @@ # automatic memory management a = nothing

Beyond memory management, there are a whole range of array operations to process your data. This includes several higher-order operations that take other code as arguments, such as map, reduce or broadcast. With these, it is possible to perform kernel-like operations without actually writing your own GPU kernels:

a = Metal.zeros(1024)
 b = Metal.ones(1024)
-a.^2 .+ sin.(b)
+a.^2 .+ sin.(b)