From 5ccad5b9a0007322392bf8af2d72dac0bae321e6 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 22 Nov 2025 22:28:50 -0400
Subject: [PATCH] Simdgroup intrinsic attributes and tweak llvm triple

---
 src/metal.jl | 12 ++++++++++--
 1 file changed, 10 insertions(+), 2 deletions(-)

diff --git a/src/metal.jl b/src/metal.jl
index d3a83d61..24ae1e6c 100644
--- a/src/metal.jl
+++ b/src/metal.jl
@@ -26,7 +26,7 @@ source_code(target::MetalCompilerTarget) = "text"
 # Metal is not supported by our LLVM builds, so we can't get a target machine
 llvm_machine(::MetalCompilerTarget) = nothing
 
-llvm_triple(target::MetalCompilerTarget) = "air64-apple-macosx$(target.macos)"
+llvm_triple(target::MetalCompilerTarget) = "air64_v$(target.air.major)$(target.air.minor)-apple-macosx$(target.macos)"
 
 llvm_datalayout(target::MetalCompilerTarget) =
     "e-p:64:64:64"*
@@ -1018,7 +1018,7 @@ function annotate_air_intrinsics!(@nospecialize(job::CompilerJob), mod::LLVM.Mod
 
         # synchronization
         if fn == "air.wg.barrier" || fn == "air.simdgroup.barrier"
-            add_attributes("nounwind", "convergent")
+            add_attributes("nounwind", "mustprogress", "convergent", "willreturn")
 
         # atomics
         elseif match(r"air.atomic.(local|global).load", fn) !== nothing
@@ -1033,6 +1033,14 @@ function annotate_air_intrinsics!(@nospecialize(job::CompilerJob), mod::LLVM.Mod
         elseif match(r"^air.atomic.(local|global).(add|sub|min|max|and|or|xor)", fn) !== nothing
             # TODO: "memory(argmem: readwrite)" on LLVM 16+
             add_attributes("argmemonly", "nounwind")
+
+        # simdgroup
+        elseif match(r"air.simdgroup_matrix_8x8_multiply_accumulate", fn) !== nothing
+            add_attributes("convergent", "mustprogress", "nounwind", "willreturn")
+        elseif match(r"air.simdgroup_matrix_8x8_load", fn) !== nothing
+            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")
         end
     end