Metal: add_parameter_address_spaces! creates a phi with mismatched address spaces #1004

On Metal, a kernel that selects between an array element and a scalar argument fails to compile:

using Metal

function kernel(out, x, neutral, n)
    i = thread_position_in_grid().x
    v = i <= n ? x[i] : neutral   # without `@inbounds`
    n > 256 && return
    @inbounds out[i] = v
    return
end

@metal threads=64 kernel(MtlArray(zeros(Int32, 64)), MtlArray(Int32.(1:32)), Int32(-1), 32)

ERROR: Compilation to AIR failed: Failed to downgrade bitcode to LLVM 14.0.0:
<input>: error: Invalid phi record (Producer: 'LLVM18.1.7jl' Reader: 'LLVM 23.1.1jl')

The bitcode is already invalid before the downgrader reads it (llvm-as on the dumped .ll: '%2' defined with type 'ptr' but expected 'ptr addrspace(1)'). Verifying after each step of finish_ir! shows it breaks in add_parameter_address_spaces!:

1. pass_by_reference! turns neutral::Int32 into a ptr addrspace(1) parameter.
2. Optimization turns the select into a load from phi ptr addrspace(1) [ %x_i, %L131 ], [ %neutral, %entry ].
3. add_parameter_address_spaces! treats that parameter as byref (src/metal.jl:1516), copies it into an AS0 alloca, and maps the AS1 parameter to that AS0 pointer. Loads don't care, but the phi now has an AS0 incoming value.

Skipping the copy for by-ref parameters that are already in a non-generic address space fixes it:

-            elseif byref[i]
+            elseif byref[i] && param.addrspace == 0

With that change (tested by overriding the method), the reproducer compiles and gives the right result, and so do the KernelAbstractions @groupreduce kernels (JuliaGPU/KernelAbstractions.jl#830) that originally hit this. With @inbounds x[i] the phi only forms after the pass, which is why this is rare.

It would also help to run LLVM.verify after finish_ir!, at least in tests, so that such IR fails at its source rather than in the downgrader.

GPUCompiler 2.13.2 (the code is unchanged on main), LLVMDowngrader_jll 0.12.0, Julia 1.12.7, M1.