Metal programming in Julia
With Metal.jl it's possible to program GPUs on macOS using the Metal programming framework.
The package is a work-in-progress. There are bugs, functionality is missing, and performance hasn't been optimized. Expect to have to make changes to this package if you want to use it. PRs are very welcome!
These requirements are fairly strict, and are due to our limited development resources (manpower, hardware). Supported OS versions follow Apple's support, which is typically the three latest major verions.
Metal.jl can be installed with the Julia package manager. From the Julia REPL, type ] to enter the Pkg REPL mode and run:
Or, equivalently, via the Pkg API:
julia> import Pkg; Pkg.add("Metal")
For an overview of the toolchain in use, you can run the following command after importing the package:
julia> using Metal julia> Metal.versioninfo() macOS 15.5.0, Darwin 24.5.0 Toolchain: - Julia: 1.11.6 - LLVM: 16.0.6 Julia packages: - Metal.jl: 1.6.3 - GPUArrays: 11.2.3 - GPUCompiler: 1.6.1 - KernelAbstractions: 0.9.37 - ObjectiveC: 3.4.2 - LLVM: 9.4.2 - LLVMDowngrader_jll: 0.6.0+1 1 device: - Apple M2 Max (64.000 KiB allocated)
The easiest way to work with Metal.jl, is by using its array abstraction. The MtlArray type is both meant to be a convenient container for device memory, as well as provide a data-parallel abstraction for using the GPU without writing your own kernels:
julia> a = MtlArray([1])
1-element MtlVector{Int64, Metal.PrivateStorage}:
1
julia> a .+ 1
1-element MtlVector{Int64, Metal.PrivateStorage}:
2
The above array abstractions are all implemented using Metal kernels written in Julia. These kernels follow a similar programming style to Julia's other GPU back-ends, and with that deviate from how kernels are implemented in Metal C (i.e., indexing intrinsics are functions not arguments, arbitrary aggregate arguments are supported, etc):
julia> function vadd(a, b, c)
i = thread_position_in_grid_1d()
c[i] = a[i] + b[i]
return
end
vadd (generic function with 1 method)
julia> a = MtlArray([1,1,1,1]); b = MtlArray([2,2,2,2]); c = similar(a);
julia> @metal threads=2 groups=2 vadd(a, b, c);
julia> Array(c)
4-element Vector{Int64}:
3
3
3
3
Finally, all of the above functionality is made possible by interfacing with the Metal libraries through ObjectiveC.jl. We provide low-level objects and functions that map These low-level API wrappers, along with some slightly higher-level Julia wrappers, are available in the MTL submodule exported by Metal.jl:
julia> using Metal
julia> dev = device()
<AGXG13XDevice: 0x14c17f200>
name = Apple M1 Pro
julia> dev.name
NSString("Apple M1 Pro")
This package builds upon the experience of several Julia contributors to CUDA.jl, AMDGPU.jl and oneAPI.jl.
RetroSearch is an open source project built by @garambo | Open a GitHub Issue
Search and Browse the WWW like it's 1997 | Search results from DuckDuckGo
HTML:
3.2
| Encoding:
UTF-8
| Version:
0.7.4