[maleadt]

KernelAbstractions 0.10 moves its back-end API into a separate package, KernelInterface. A back end implements KernelInterface: allocating and copying memory, selecting devices, compiling and launching kernels, and the device-side intrinsics. KernelAbstractions then launches @kernel kernels the same way on every back end (JuliaGPU/KernelAbstractions.jl#801). This PR ports Metal.jl to that model in one step. It replaces #917, which added KernelInterface next to the existing KernelAbstractions 0.9 back end, and the kaintr branch, which ported an earlier version of it to KernelAbstractions 0.10.

MetalBackend now implements KernelInterface, and Metal depends on KernelInterface instead of KernelAbstractions. That makes the KernelInterface layer usable on its own, without KernelAbstractions' macros:

using Metal
import KernelInterface as KI

function vadd(c, a, b)
    i = KI.get_global_id().x
    if i <= length(c)
        @inbounds c[i] = a[i] + b[i]
    end
    return
end

a = Metal.rand(Float32, 1000); b = Metal.rand(Float32, 1000); c = similar(a)
KI.@launch MetalBackend() ndrange=length(c) vadd(c, a, b)

KernelAbstractions becomes a weak dependency. Its extension only provides what KernelAbstractions still needs from a back end: the stack allocation behind @private, and the rule for moving an MtlArray to the CPU with adapt. Metal's own copy of the KernelAbstractions launch path goes away: partitioning the ndrange, building the kernel's context, sizing the threadgroups, calling the kernel, and the overrides of KernelAbstractions' index functions. KI.copyto! now also accepts contiguous views of host arrays, and adapt(MetalBackend(), x) moves any array to the GPU, as adapt(MtlArray, x) does, not only Arrays. Threadgroups are still as large as the pipeline's maxTotalThreadsPerThreadgroup allows, which is what KernelInterface's default launch_configuration picks, so Metal doesn't override it.

Kernels behave the same, but they are launched differently. KernelAbstractions now launches them on a 3-D grid, and computes @index in 32 bits when the iteration space fits (JuliaGPU/KernelAbstractions.jl#797). So a kernel over a 3-D ndrange doesn't need divisions to compute its index, which the test suite checks for the generated LLVM IR. On Metal this matters a lot, since 64-bit integer division is slow on Apple GPUs. With a KA kernel that copies a Float32 array, on an M1 (GPU time of the command buffer):

	main (KA 0.9)	this PR
256×256×256, @index(Global, Cartesian)	12.5 ms	2.3 ms
256×256×256, @index(Global, Linear)	11.4 ms	2.4 ms
255×257×129, @index(Global, Cartesian)	6.4 ms	1.1 ms
255×257×129, @index(Global, Linear)	5.9 ms	1.1 ms
launch of an empty kernel	2.33 µs, 768 bytes	2.58 µs, 768 bytes
launch of a kernel with 28 arguments	9.6 µs, 4480 bytes	3.1 µs, 2112 bytes

The kernel times vary by about 10% between runs of this PR. The launch benchmark uses 28 arguments rather than the 40 of CUDA.jl's, because a Metal kernel takes at most 31 buffers, and every argument, the kernel's context and Metal's kernel state each take one.

KernelInterface 0.4 passes the arguments to the back end as a tuple (JuliaGPU/KernelAbstractions.jl#811). Metal has no public launch function that takes a tuple, but calling a HostKernel forwards its arguments as one to an internal function, and KI.launch calls that function directly, so nothing is splatted. KI.launch rejects threads and groups, which would override the launch geometry KernelInterface has validated, but passes Metal's other launch options on:

kernel = KI.@launch MetalBackend() launch=false vadd(c, a, b)
kernel(c, a, b; ndrange=length(c), queue=global_queue(device()))   # fine
kernel(c, a, b; ndrange=length(c), threads=256)                     # ArgumentError

The first commit is the port. The next ones implement parts of KernelInterface that are optional: sub-groups (SIMD-groups, including a partial last one), ordering work across tasks with Metal events for KernelAbstractions.@spawn, where the waiting task's command buffer waits on the GPU instead of blocking the host, and KI.versioninfo.

KernelAbstractions' conversion tests, skipped until now, pass with 32-bit indices. Its test that launches a kernel with 41 arguments is skipped, since a Metal kernel takes at most 31 buffers, one per argument. That limit isn't new, but it now shows in KernelAbstractions' test suite. Unlike CUDA.jl, Metal already converted the arguments twice (once for the types to compile for, and again while encoding them), so KernelAbstractions' generic launch doesn't add a conversion.

This needs a breaking release, since it requires KernelAbstractions 0.10, which isn't registered yet. Until then, the last commit takes KernelAbstractions and KernelInterface from their development branch through [sources]. Julia 1.10 ignores [sources] and Julia 1.11 doesn't support workspaces, so on Buildkite those two set up the test environment by hand and run the tests in it directly. That commit is dropped before merging.

[github-actions]

Metal Benchmarks

Details

Benchmark suite	Current: 942b315	Previous: 76951f2	Ratio
array/accumulate/Float32/1d	393583 ns	391959 ns	1.00
array/accumulate/Float32/dims=1	369250 ns	370167 ns	1.00
array/accumulate/Float32/dims=1L	8864875 ns	8878291 ns	1.00
array/accumulate/Float32/dims=2	433167 ns	427250 ns	1.01
array/accumulate/Float32/dims=2L	2803000 ns	2669209 ns	1.05
array/accumulate/Int64/1d	845125 ns	849083 ns	1.00
array/accumulate/Int64/dims=1	909584 ns	913084 ns	1.00
array/accumulate/Int64/dims=1L	9551208 ns	9560458 ns	1.00
array/accumulate/Int64/dims=2	1218125 ns	1210792 ns	1.01
array/accumulate/Int64/dims=2L	6563833 ns	6543625 ns	1.00
array/broadcast	241709 ns	224916 ns	1.07
array/construct	2250 ns	2291 ns	0.98
array/permutedims/2d	341709 ns	455458 ns	0.75
array/permutedims/3d	464042 ns	1028542 ns	0.45
array/permutedims/4d	625041 ns	1119334 ns	0.56
array/private/copy	224958 ns	228583 ns	0.98
array/private/copyto!/cpu_to_gpu	229000 ns	213417 ns	1.07
array/private/copyto!/gpu_to_cpu	217958 ns	215708 ns	1.01
array/private/copyto!/gpu_to_gpu	214584 ns	221916 ns	0.97
array/private/iteration/findall/bool	1047583 ns	1051292 ns	1.00
array/private/iteration/findall/int	1215250 ns	1217584 ns	1.00
array/private/iteration/findfirst/bool	1011125 ns	1008458 ns	1.00
array/private/iteration/findfirst/int	1019709 ns	1016083 ns	1.00
array/private/iteration/findmin/1d	1087792 ns	1073084 ns	1.01
array/private/iteration/findmin/2d	1024875 ns	1020750 ns	1.00
array/private/iteration/logical	1539833 ns	1551500 ns	0.99
array/private/iteration/scalar	1381167 ns	1398250 ns	0.99
array/random/rand/Float32	444209 ns	407458 ns	1.09
array/random/rand/Int64	521208 ns	499458 ns	1.04
array/random/rand!/Float32	424583 ns	401125 ns	1.06
array/random/rand!/Int64	443417 ns	426625 ns	1.04
array/random/randn/Float32	369458 ns	391666 ns	0.94
array/random/randn!/Float32	355833 ns	375875 ns	0.95
array/reductions/mapreduce/Float32/1d	270042 ns	277875 ns	0.97
array/reductions/mapreduce/Float32/dims=1	353875 ns	353792 ns	1.00
array/reductions/mapreduce/Float32/dims=1L	612959 ns	603417 ns	1.02
array/reductions/mapreduce/Float32/dims=2	356375 ns	358417 ns	0.99
array/reductions/mapreduce/Float32/dims=2L	1220125 ns	1236416 ns	0.99
array/reductions/mapreduce/Int64/1d	459208 ns	459958 ns	1.00
array/reductions/mapreduce/Int64/dims=1	635584 ns	634125 ns	1.00
array/reductions/mapreduce/Int64/dims=1L	1005583 ns	1014250 ns	0.99
array/reductions/mapreduce/Int64/dims=2	777375 ns	788625 ns	0.99
array/reductions/mapreduce/Int64/dims=2L	2205917 ns	2200083 ns	1.00
array/reductions/reduce/Float32/1d	271583 ns	271167 ns	1.00
array/reductions/reduce/Float32/dims=1	355958 ns	359625 ns	0.99
array/reductions/reduce/Float32/dims=1L	615042 ns	618458 ns	0.99
array/reductions/reduce/Float32/dims=2	245708 ns	247334 ns	0.99
array/reductions/reduce/Float32/dims=2L	468500 ns	463500 ns	1.01
array/reductions/reduce/Int64/1d	455166 ns	462250 ns	0.98
array/reductions/reduce/Int64/dims=1	639000 ns	639250 ns	1.00
array/reductions/reduce/Int64/dims=1L	1007833 ns	1007959 ns	1.00
array/reductions/reduce/Int64/dims=2	257083 ns	258125 ns	1.00
array/reductions/reduce/Int64/dims=2L	656667 ns	663458 ns	0.99
array/shared/copy	129959 ns	128000 ns	1.02
array/shared/copyto!/cpu_to_gpu	37500 ns	38292 ns	0.98
array/shared/copyto!/gpu_to_cpu	38542 ns	36708 ns	1.05
array/shared/copyto!/gpu_to_gpu	39000 ns	36792 ns	1.06
array/shared/iteration/findall/bool	1052375 ns	1061167 ns	0.99
array/shared/iteration/findall/int	1211958 ns	1223625 ns	0.99
array/shared/iteration/findfirst/bool	974834 ns	932917 ns	1.04
array/shared/iteration/findfirst/int	947584 ns	987583 ns	0.96
array/shared/iteration/findmin/1d	1045167 ns	1055750 ns	0.99
array/shared/iteration/findmin/2d	1029708 ns	1028333 ns	1.00
array/shared/iteration/logical	1521167 ns	1512792 ns	1.01
array/shared/iteration/scalar	76.70328542094455 ns	76.42374616171955 ns	1.00
array/sorting/1d	1910875 ns	1912167 ns	1.00
array/sorting/2d	8344500 ns	8268833 ns	1.01
integration/byval/reference	1118959 ns	1122583 ns	1.00
integration/byval/slices=1	1124792 ns	1119417 ns	1.00
integration/byval/slices=2	2032750 ns	2029125 ns	1.00
integration/byval/slices=3	6798916 ns	6760792 ns	1.01
integration/metaldevrt	394292 ns	397583 ns	0.99
kernel/indexing	212750 ns	212542 ns	1.00
kernel/indexing_checked	389875 ns	396666 ns	0.98
kernel/launch	2120.3333333333335 ns	2092.5555555555557 ns	1.01
kernel/rand	407042 ns	408917 ns	1.00
latency/import	1732971750 ns	1858464250 ns	0.93
latency/precompile	25785776708 ns	31532892250 ns	0.82
latency/ttfp	2222702541 ns	2289033959 ns	0.97
metal/synchronization/context	535.5263157894736 ns	530.921052631579 ns	1.01
metal/synchronization/stream	628.4093567251462 ns	618.2151162790698 ns	1.02

This comment was automatically generated by workflow using github-action-benchmark.

[christiangnrd]

This assumes the work group is always the max which doesn't feel right. I think it would be more useful to have it take into account the work groups size. (I can open an issue in KA if that would be better)

[github-actions]

Metal.jl Benchmarks

Details

Benchmark suite	Current: 97e6dbd	Previous: b8d46ce	Ratio
array/accumulate/Float32/1d	394167 ns	579959 ns	0.68
array/accumulate/Float32/dims=1	370458 ns	571958 ns	0.65
array/accumulate/Float32/dims=1L	8868791 ns	8806687.5 ns	1.01
array/accumulate/Float32/dims=2	429291 ns	621416.5 ns	0.69
array/accumulate/Float32/dims=2L	2587708 ns	3024708 ns	0.86
array/accumulate/Int64/1d	773333 ns	902083 ns	0.86
array/accumulate/Int64/dims=1	875292 ns	1112334 ns	0.79
array/accumulate/Int64/dims=1L	3548375 ns	9482125 ns	0.37
array/accumulate/Int64/dims=2	1162292 ns	1433583 ns	0.81
array/accumulate/Int64/dims=2L	16099333 ns	6482270.5 ns	2.48
array/broadcast	240209 ns	371084 ns	0.65
array/construct	2250 ns	3666.5 ns	0.61
array/permutedims/2d	340459 ns	730125 ns	0.47
array/permutedims/3d	458166 ns	1255500 ns	0.36
array/permutedims/4d	616208 ns	1677792 ns	0.37
array/private/copy	222667 ns	402625 ns	0.55
array/private/copyto!/cpu_to_gpu	218250 ns	394875 ns	0.55
array/private/copyto!/gpu_to_cpu	222834 ns	393458 ns	0.57
array/private/copyto!/gpu_to_gpu	218375 ns	362375 ns	0.60
array/private/iteration/findall/bool	529417 ns	1122520.5 ns	0.47
array/private/iteration/findall/int	608083 ns	1285709 ns	0.47
array/private/iteration/findfirst/bool	374833 ns	1444000 ns	0.26
array/private/iteration/findfirst/int	378209 ns	1461083.5 ns	0.26
array/private/iteration/findmin/1d	306625 ns	1556583.5 ns	0.20
array/private/iteration/findmin/2d	1106833 ns	1248916 ns	0.89
array/private/iteration/logical	544459 ns	1814937.5 ns	0.30
array/private/iteration/scalar	1391792 ns	2995875 ns	0.46
array/random/rand/Float32	440958 ns	606541 ns	0.73
array/random/rand/Int64	514792 ns	606250 ns	0.85
array/random/rand!/Float32	425042 ns	534458 ns	0.80
array/random/rand!/Int64	452416 ns	505792 ns	0.89
array/random/randn/Float32	369417 ns	564125 ns	0.65
array/random/randn!/Float32	354250 ns	490500 ns	0.72
array/reductions/mapreduce/Float32/1d	224750 ns	731084 ns	0.31
array/reductions/mapreduce/Float32/dims=1	388792 ns	491187.5 ns	0.79
array/reductions/mapreduce/Float32/dims=1L	581458 ns	678895.5 ns	0.86
array/reductions/mapreduce/Float32/dims=2	316250 ns	508292 ns	0.62
array/reductions/mapreduce/Float32/dims=2L	454958 ns	1349166 ns	0.34
array/reductions/mapreduce/Int64/1d	239125 ns	920291 ns	0.26
array/reductions/mapreduce/Int64/dims=1	401209 ns	789042 ns	0.51
array/reductions/mapreduce/Int64/dims=1L	1005250 ns	1065667 ns	0.94
array/reductions/mapreduce/Int64/dims=2	331417 ns	945729.5 ns	0.35
array/reductions/mapreduce/Int64/dims=2L	607834 ns	2226375 ns	0.27
array/reductions/reduce/Float32/1d	230541 ns	731333.5 ns	0.32
array/reductions/reduce/Float32/dims=1	389792 ns	494041.5 ns	0.79
array/reductions/reduce/Float32/dims=1L	585084 ns	680270.5 ns	0.86
array/reductions/reduce/Float32/dims=2	319959 ns	405500 ns	0.79
array/reductions/reduce/Float32/dims=2L	417709 ns	498250 ns	0.84
array/reductions/reduce/Int64/1d	251375 ns	928604.5 ns	0.27
array/reductions/reduce/Int64/dims=1	327292 ns	793875 ns	0.41
array/reductions/reduce/Int64/dims=1L	1002750 ns	1063583 ns	0.94
array/reductions/reduce/Int64/dims=2	330791 ns	322500 ns	1.03
array/reductions/reduce/Int64/dims=2L	622459 ns	695208 ns	0.90
array/shared/copy	128875 ns	146375 ns	0.88
array/shared/copyto!/cpu_to_gpu	37125 ns	40750 ns	0.91
array/shared/copyto!/gpu_to_cpu	37875 ns	38708 ns	0.98
array/shared/copyto!/gpu_to_gpu	38208 ns	39292 ns	0.97
array/shared/iteration/findall/bool	530459 ns	1129083 ns	0.47
array/shared/iteration/findall/int	605834 ns	1293458 ns	0.47
array/shared/iteration/findfirst/bool	378541 ns	1145583 ns	0.33
array/shared/iteration/findfirst/int	377750 ns	1165042 ns	0.32
array/shared/iteration/findmin/1d	372750 ns	1285958 ns	0.29
array/shared/iteration/findmin/2d	1106750 ns	1250458 ns	0.89
array/shared/iteration/logical	536791 ns	1642583 ns	0.33
array/shared/iteration/scalar	85.17695473251028 ns	3729.125 ns	0.022841002844503813
array/sorting/1d	1920750 ns	2336750 ns	0.82
array/sorting/2d	8329542 ns	8350375 ns	1.00
integration/byval/reference	1119583 ns	1144000 ns	0.98
integration/byval/slices=1	1104625 ns	1138875 ns	0.97
integration/byval/slices=2	2023875 ns	2042917 ns	0.99
integration/byval/slices=3	6851750 ns	7087083 ns	0.97
integration/metaldevrt	393334 ns	446042 ns	0.88
kernel/indexing	214416 ns	361000 ns	0.59
kernel/indexing_checked	392792 ns	545000 ns	0.72
kernel/launch	1529.1 ns	2062.5 ns	0.74
kernel/rand	407333 ns	534958 ns	0.76
latency/import	2006394209 ns	1999970250 ns	1.00
latency/precompile	30904116292 ns	36823452646 ns	0.84
latency/ttfp	2243051250 ns	2211888562.5 ns	1.01
metal/synchronization/context	530.921052631579 ns	541.8854166666666 ns	0.98
metal/synchronization/stream	654.8395061728395 ns	274.2063492063492 ns	2.39

This comment was automatically generated by workflow using github-action-benchmark.

[codecov]

Codecov Report

❌ Patch coverage is 99.18033% with 1 line in your changes missing coverage. Please review.
✅ Project coverage is 86.82%. Comparing base (8c421b8) to head (97e6dbd).
⚠️ Report is 2 commits behind head on main.

Files with missing lines	Patch %	Lines
src/MetalKernels.jl	99.12%	1 Missing ⚠️

Additional details and impacted files

@@            Coverage Diff             @@
##             main     #994      +/-   ##
==========================================
+ Coverage   86.51%   86.82%   +0.30%     
==========================================
  Files          89       89              
  Lines        6326     6419      +93     
==========================================
+ Hits         5473     5573     +100     
+ Misses        853      846       -7     

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:

• ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

[yolhan83]

Hello, just while testing :

julia> using Atomix, Metal, KernelAbstractions

julia> @kernel function atomictest(out)
           Atomix.@atomic out[1] += 1f0
       end

julia> out = MtlArray(zeros(Float32, 1))
1-element MtlVector{Float32, Metal.SharedStorage}:
 0.0

julia> ker = atomictest(MetalBackend(), 32)
KernelAbstractions.Kernel{MetalBackend, KernelAbstractions.NDIteration.StaticSize{(32,)}, KernelAbstractions.NDIteration.DynamicSize, typeof(gpu_atomictest)}(MetalBackend(), gpu_atomictest)

julia> ker(out; ndrange=32)
ERROR: Compilation to native code failed; see below for details.
If you think this is a bug, please file an issue and attach:
- /var/folders/px/4l1mqd354pbdzzdp94rgc40h0000gn/T/jl_DiNqzCjQdA.air
- /var/folders/px/4l1mqd354pbdzzdp94rgc40h0000gn/T/jl_DiNqzCjQdA.metallib
Stacktrace:
  [1] error(s::String)
    @ Base ./error.jl:44
  [2] macro expansion
    @ ~/.julia/packages/Metal/uNxGY/src/compiler/compilation.jl:567 [inlined]
  [3] macro expansion
    @ ~/.julia/packages/ObjectiveC/lG1iL/src/os.jl:264 [inlined]
  [4] macro expansion
    @ ~/.julia/packages/Metal/uNxGY/src/compiler/compilation.jl:556 [inlined]
  [5] (::Metal.var"#link_pipeline##0#link_pipeline##1"{Metal.MTL.MTLDevice, Vector{UInt8}, Vector{UInt8}, String})()
    @ Metal ~/.julia/packages/ObjectiveC/lG1iL/src/foundation.jl:733
  [6] macro expansion
    @ ~/.julia/packages/ObjectiveC/lG1iL/src/foundation.jl:661 [inlined]
  [7] macro expansion
    @ ./lock.jl:376 [inlined]
  [8] ObjectiveC.Foundation.NSAutoreleasePool(f::Metal.var"#link_pipeline##0#link_pipeline##1"{…})
    @ ObjectiveC.Foundation ~/.julia/packages/ObjectiveC/lG1iL/src/foundation.jl:653
  [9] link_pipeline
    @ ~/.julia/packages/Metal/uNxGY/src/compiler/compilation.jl:552 [inlined]
 [10] macro expansion
    @ ~/.julia/packages/Metal/uNxGY/src/compiler/execution.jl:254 [inlined]
 [11] macro expansion
    @ ./lock.jl:376 [inlined]
 [12] link_kernel(mi::Core.MethodInstance; name::Nothing, kwargs::@Kwargs{})
    @ Metal ~/.julia/packages/Metal/uNxGY/src/compiler/execution.jl:237
 [13] mtlfunction(f::typeof(gpu_atomictest), tt::Type{…}; source::typeof(gpu_atomictest), kwargs::@Kwargs{…})
    @ Metal ~/.julia/packages/Metal/uNxGY/src/compiler/execution.jl:231
 [14] mtlfunction
    @ ~/.julia/packages/Metal/uNxGY/src/compiler/execution.jl:230 [inlined]
 [15] #kernel_function#2
    @ ~/.julia/packages/Metal/uNxGY/src/MetalKernels.jl:102 [inlined]
 [16] kernel_function
    @ ~/.julia/packages/Metal/uNxGY/src/MetalKernels.jl:99 [inlined]
 [17] launch_kernel(obj::KernelAbstractions.Kernel{…}, launch::KernelAbstractions.NDLaunch{…}, subgroups::Nothing, ndrange::Tuple{…}, _workgroupsize::Nothing, iterspace::KernelAbstractions.NDIteration.NDRange{…}, args::Tuple{…})
    @ KernelAbstractions ~/.julia/packages/KernelAbstractions/tHp3q/src/backend_launch.jl:132
 [18] launch_kernel
    @ ~/.julia/packages/KernelAbstractions/tHp3q/src/backend_launch.jl:122 [inlined]
 [19] launch_tuple(obj::KernelAbstractions.Kernel{…}, args::Tuple{…}; ndrange::Int64, workgroupsize::Nothing)
    @ KernelAbstractions ~/.julia/packages/KernelAbstractions/tHp3q/src/backend_launch.jl:109
 [20] launch_tuple
    @ ~/.julia/packages/KernelAbstractions/tHp3q/src/backend_launch.jl:98 [inlined]
 [21] top-level scope
    @ REPL[32]:1
 [22] top-level scope
    @ ~/.julia/packages/Metal/uNxGY/src/initialization.jl:106

caused by: NSError: Compilation failed due to an interrupted connection: XPC_ERROR_CONNECTION_INTERRUPTED. This error occurred after multiple retries. (AGXMetalG17X, code 2)
Stacktrace:
  [1] throw_error(err::ObjectiveC.id{ObjectiveC.Foundation.NSError})
    @ Metal.MTL ~/.julia/packages/Metal/uNxGY/lib/mtl/MTL.jl:23
  [2] MTLComputePipelineState
    @ ~/.julia/packages/Metal/uNxGY/lib/mtl/compute_pipeline.jl:25 [inlined]
  [3] macro expansion
    @ ~/.julia/packages/Metal/uNxGY/src/compiler/compilation.jl:560 [inlined]
  [4] macro expansion
    @ ~/.julia/packages/ObjectiveC/lG1iL/src/os.jl:264 [inlined]
  [5] macro expansion
    @ ~/.julia/packages/Metal/uNxGY/src/compiler/compilation.jl:556 [inlined]
  [6] (::Metal.var"#link_pipeline##0#link_pipeline##1"{Metal.MTL.MTLDevice, Vector{UInt8}, Vector{UInt8}, String})()
    @ Metal ~/.julia/packages/ObjectiveC/lG1iL/src/foundation.jl:733
  [7] macro expansion
    @ ~/.julia/packages/ObjectiveC/lG1iL/src/foundation.jl:661 [inlined]
  [8] macro expansion
    @ ./lock.jl:376 [inlined]
  [9] ObjectiveC.Foundation.NSAutoreleasePool(f::Metal.var"#link_pipeline##0#link_pipeline##1"{…})
    @ ObjectiveC.Foundation ~/.julia/packages/ObjectiveC/lG1iL/src/foundation.jl:653
 [10] link_pipeline
    @ ~/.julia/packages/Metal/uNxGY/src/compiler/compilation.jl:552 [inlined]
 [11] macro expansion
    @ ~/.julia/packages/Metal/uNxGY/src/compiler/execution.jl:254 [inlined]
 [12] macro expansion
    @ ./lock.jl:376 [inlined]
 [13] link_kernel(mi::Core.MethodInstance; name::Nothing, kwargs::@Kwargs{})
    @ Metal ~/.julia/packages/Metal/uNxGY/src/compiler/execution.jl:237
 [14] mtlfunction(f::typeof(gpu_atomictest), tt::Type{…}; source::typeof(gpu_atomictest), kwargs::@Kwargs{…})
    @ Metal ~/.julia/packages/Metal/uNxGY/src/compiler/execution.jl:231
 [15] mtlfunction
    @ ~/.julia/packages/Metal/uNxGY/src/compiler/execution.jl:230 [inlined]
 [16] #kernel_function#2
    @ ~/.julia/packages/Metal/uNxGY/src/MetalKernels.jl:102 [inlined]
 [17] kernel_function
    @ ~/.julia/packages/Metal/uNxGY/src/MetalKernels.jl:99 [inlined]
 [18] launch_kernel(obj::KernelAbstractions.Kernel{…}, launch::KernelAbstractions.NDLaunch{…}, subgroups::Nothing, ndrange::Tuple{…}, _workgroupsize::Nothing, iterspace::KernelAbstractions.NDIteration.NDRange{…}, args::Tuple{…})
    @ KernelAbstractions ~/.julia/packages/KernelAbstractions/tHp3q/src/backend_launch.jl:132
 [19] launch_kernel
    @ ~/.julia/packages/KernelAbstractions/tHp3q/src/backend_launch.jl:122 [inlined]
 [20] launch_tuple(obj::KernelAbstractions.Kernel{…}, args::Tuple{…}; ndrange::Int64, workgroupsize::Nothing)
    @ KernelAbstractions ~/.julia/packages/KernelAbstractions/tHp3q/src/backend_launch.jl:109
 [21] launch_tuple
    @ ~/.julia/packages/KernelAbstractions/tHp3q/src/backend_launch.jl:98 [inlined]
 [22] top-level scope
    @ REPL[32]:1
 [23] top-level scope
    @ ~/.julia/packages/Metal/uNxGY/src/initialization.jl:106
Some type information was truncated. Use `show(err)` to see complete types.

it works with Metal.@atomic though.

[maleadt]

I'll take a look. In the future, please open an issue; this problem is unrelated to this PR.