Legalize barriers for the target AIR version here, not in GPUCompiler

src/compiler/compilation.jl

@@ -237,8 +237,7 @@ function GPUCompiler.finish_ir!(@nospecialize(job::MetalCompilerJob),
 
     # device code emits intrinsics in their newest form; rewrite them for older AIR versions.
     # this lives here, next to the emitters, so that the two change together. GPUCompiler
-    # only legalizes what it emits itself (atomics and fences), and barriers, which share
-    # their ABI.
+    # only legalizes the intrinsics it emits itself (atomics and fences).
     ## simdgroup
     if job.config.target.air < v"2.8"
         # AIR 2.8 generalized the simdgroup matrix load/store intrinsics, replacing
@@ -352,6 +351,37 @@ function GPUCompiler.finish_ir!(@nospecialize(job::MetalCompilerJob),
             erase!(f)
         end
     end
+    ## barriers
+    # AIR 2.9 added a memory order operand between the memory flags and the scope, which
+    # before MSL 4.1 has to be sequentially consistent
+    for fn in ("air.wg.barrier", "air.simdgroup.barrier")
+        f = get(mod.functions, fn, nothing)
+        f === nothing && continue
+        T_i32 = LLVM.Int32Type()
+        if job.config.target.air >= v"2.9"
+            job.config.target.metal >= v"4.1" && continue
+            for call in collect(f.users)
+                call::LLVM.CallInst
+                call.operands[2] = ConstantInt(T_i32, 5)
+            end
+            continue
+        end
+
+        old_ft = LLVM.FunctionType(LLVM.VoidType(), [T_i32, T_i32])
+        f.name = fn * ".air29"
+        old_f = LLVM.Function(mod, fn, old_ft)
+        append!(old_f.function_attributes, f.function_attributes)
+        for call in collect(f.users)
+            call::LLVM.CallInst
+            @dispose builder=IRBuilder() begin
+                position!(builder, LLVM.before(call))
+                flags, _, scope = collect(call.arguments)
+                call!(builder, old_ft, old_f, [flags, scope])
+                erase!(call)
+            end
+        end
+        erase!(f)
+    end
 
     # pointer type information for typed intrinsics
     # (this is consumed by the LLVM IR downgrader)

src/device/intrinsics/synchronization.jl

@@ -94,7 +94,7 @@ end
 @device_function @inline simdgroup_barrier(flags, scope::thread_scope) =
     simdgroup_barrier(flags, memory_order_seq_cst, scope)
 
-# barriers are emitted in the form of AIR 2.9, with a memory order, which GPUCompiler drops
+# barriers are emitted in the form of AIR 2.9, with a memory order, which `finish_ir!` drops
 # for older targets (and strengthens to sequentially consistent before MSL 4.1)
 @device_function @inline function barrier(::Val{group}, ::Val{flags}, ::Val{order},
                                           ::Val{scope}) where {group, flags, order, scope}

test/device/intrinsics/synchronization.jl

@@ -78,14 +78,14 @@
             simdgroup_barrier(Metal.MemoryFlagThreadGroup)
             return
         end
-        # the AIR 2.9 form, with an order, which GPUCompiler legalizes for the target
+        # the AIR 2.9 form, with an order, legalized for the target
         @testset "Metal $metal, AIR $air" for (metal, air) in ((v"3.2", v"2.7"),
                                                               (v"4.0", v"2.9"),
                                                               (v"4.1", v"2.9"))
             ir = sprint(io -> Metal.code_air(io, barrier_abi,
                                              Tuple{Core.LLVMPtr{Int32,Metal.AS.Device}};
                                              kernel=true, metal, air))
-            # (before MSL 4.1, GPUCompiler makes barriers sequentially consistent)
+            # (before MSL 4.1, barriers are sequentially consistent)
             release, acquire = metal >= v"4.1" ? (3, 2) : (5, 5)
             if air >= v"2.9"
                 @test occursin("@air.wg.barrier(i32 1, i32 $release, i32 2)", ir)