diff --git a/dev/.documenter-siteinfo.json b/dev/.documenter-siteinfo.json index 16007cc14..b4fa12e3e 100644 --- a/dev/.documenter-siteinfo.json +++ b/dev/.documenter-siteinfo.json @@ -1 +1 @@ -{"documenter":{"julia_version":"1.10.4","generation_timestamp":"2024-07-25T03:13:57","documenter_version":"1.5.0"}} \ No newline at end of file +{"documenter":{"julia_version":"1.10.4","generation_timestamp":"2024-08-06T04:33:18","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 ae77cfd5f..3e11b336e 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 46c10939b..14c4e051f 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 dac77e556..ee1c635cc 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 e93d60fb1..62ac9984c 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 40d3d7f18..bec967ecf 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 70918d6f3..40b0d076f 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 2701536d9..2e9b3ba25 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 aba3067fb..96e6fd3c3 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/profiling/index.html b/dev/profiling/index.html index 56c79da5a..66220f5b5 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/usage/array/index.html b/dev/usage/array/index.html index ed375eee5..02ae4cf2e 100644 --- a/dev/usage/array/index.html +++ b/dev/usage/array/index.html @@ -50,4 +50,4 @@ julia> Base.mapreducedim!(identity, +, b, a) 1×1 MtlMatrix{Float32, Metal.PrivateStorage}: - 6.0 + 6.0 diff --git a/dev/usage/kernel/index.html b/dev/usage/kernel/index.html index 90231ac18..7155e81d3 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 An Introduction to GPU Programming course from University of Illinois (primarily in CUDA, but the concepts are transferable)

diff --git a/dev/usage/overview/index.html b/dev/usage/overview/index.html index 5cf1f2b64..020428283 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)