[Metal] Add correct addrspace to global constants

[tgymnich]

Global constants need to be located in address space 2 in Metal:

This fixes the following use case:

using Metal
using StaticArrays

const xs = SVector(1.0f0, 2.0f0, 3.0f0)

function kernel_fma(a, b, x, out)
    i = thread_position_in_grid_1d()
    a_val = a[i]
    b_val = b[i]
    x_val = x[i]
    
    out[i] = fma(a_val, b_val, xs[x_val])

    return
end

len=1024
a = MtlArray(rand(Float32, len))
b = MtlArray(rand(Float32, len))
x = MtlArray(rand(1:3, len))
out = similar(a)

threads = 1024

@device_code_llvm dump_module=true raw=true @metal threads=threads kernel_fma(a, b, x, out)

[maleadt]

Good thing I added a test; looks like my changes trip up an LLVM assertion:

julia: /workspace/srcdir/llvm-project/llvm/lib/IR/Constants.cpp:2231: static llvm::Constant* llvm::ConstantExpr::getBitCast(llvm::Constant*, llvm::Type*, bool): Assertion `CastInst::castIsValid(Instruction::BitCast, C, DstTy) && "Invalid constantexpr bitcast!"' failed.