From be05e01e2380ce53f388094971dae8939e03d7d5 Mon Sep 17 00:00:00 2001 From: Simone Carlo Surace Date: Tue, 24 Feb 2026 16:03:47 +0100 Subject: [PATCH 1/3] Add ballot voting attributes for Metal backend --- src/metal.jl | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/src/metal.jl b/src/metal.jl index 52060ebb..400f2b6b 100644 --- a/src/metal.jl +++ b/src/metal.jl @@ -1041,6 +1041,10 @@ function annotate_air_intrinsics!(@nospecialize(job::CompilerJob), mod::LLVM.Mod add_attributes("convergent", "mustprogress", "nofree", "nounwind", "readonly", "willreturn") elseif match(r"air.simdgroup_matrix_8x8_store", fn) !== nothing add_attributes("convergent", "mustprogress", "nounwind", "willreturn", "writeonly") + + # simd voting/ballot + elseif match(r"air.simd_ballot", fn) !== nothing || match(r"air.simd_vote_all", fn) !== nothing || match(r"air.simd_vote_any", fn) !== nothing + add_attributes("convergent", "mustprogress", "nounwind", "willreturn") end end From 92ae1e43dc6d685188504d0aa34d01190a8946aa Mon Sep 17 00:00:00 2001 From: Simone Carlo Surace Date: Wed, 25 Feb 2026 09:42:11 +0100 Subject: [PATCH 2/3] Add the other simd intrinsics --- src/metal.jl | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/src/metal.jl b/src/metal.jl index 400f2b6b..321ec40a 100644 --- a/src/metal.jl +++ b/src/metal.jl @@ -1042,8 +1042,9 @@ function annotate_air_intrinsics!(@nospecialize(job::CompilerJob), mod::LLVM.Mod elseif match(r"air.simdgroup_matrix_8x8_store", fn) !== nothing add_attributes("convergent", "mustprogress", "nounwind", "willreturn", "writeonly") - # simd voting/ballot - elseif match(r"air.simd_ballot", fn) !== nothing || match(r"air.simd_vote_all", fn) !== nothing || match(r"air.simd_vote_any", fn) !== nothing + # simd intrinsics + elseif match(r"air.simd_(ballot|vote_all|vote_any|shuffle|shuffle_xor|shuffle_down|\ + shuffle_up|shuffle_and_fill_down|shuffle_and_fill_up)", fn) !== nothing add_attributes("convergent", "mustprogress", "nounwind", "willreturn") end end From b169a384fde69ef45a8bbddc01a1bcb8ed0f32a8 Mon Sep 17 00:00:00 2001 From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com> Date: Wed, 25 Feb 2026 21:00:06 -0400 Subject: [PATCH 3/3] Add even more intrinsics --- src/metal.jl | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/metal.jl b/src/metal.jl index 321ec40a..569a6cd1 100644 --- a/src/metal.jl +++ b/src/metal.jl @@ -1042,8 +1042,8 @@ function annotate_air_intrinsics!(@nospecialize(job::CompilerJob), mod::LLVM.Mod elseif match(r"air.simdgroup_matrix_8x8_store", fn) !== nothing add_attributes("convergent", "mustprogress", "nounwind", "willreturn", "writeonly") - # simd intrinsics - elseif match(r"air.simd_(ballot|vote_all|vote_any|shuffle|shuffle_xor|shuffle_down|\ + # simd permute + elseif match(r"air.simd_(ballot|all|vote_all|any|vote_any|shuffle|shuffle_xor|shuffle_down|\ shuffle_up|shuffle_and_fill_down|shuffle_and_fill_up)", fn) !== nothing add_attributes("convergent", "mustprogress", "nounwind", "willreturn") end