Updates for macOS 13 (Ventura); use bindless argument buffers

[maleadt]

If all goes well, I'll probably be bumping the requirements for Metal.jl to macOS 13 (Ventura), because it ships Metal 3 which supports bindless arguments.

[codecov]

Codecov Report

Merging #23 (e234ecd) into main (a1115cc) will decrease coverage by 1.44%.
The diff coverage is 97.56%.

@@            Coverage Diff             @@
##             main      #23      +/-   ##
==========================================
- Coverage   64.39%   62.94%   -1.45%     
==========================================
  Files          36       36              
  Lines        1070     1039      -31     
==========================================
- Hits          689      654      -35     
- Misses        381      385       +4     

Impacted Files	Coverage Δ
src/utilities.jl	75.86% <ø> (-5.22%)	⬇️
lib/core/buffer.jl	48.97% <75.00%> (+2.17%)	⬆️
lib/core/MTL.jl	90.00% <100.00%> (+40.00%)	⬆️
lib/core/device.jl	98.03% <100.00%> (+0.08%)	⬆️
src/array.jl	48.38% <100.00%> (+1.18%)	⬆️
src/compiler/execution.jl	88.17% <100.00%> (-1.60%)	⬇️
lib/core/argument_encoder.jl	51.72% <0.00%> (-27.59%)	⬇️
lib/core/command_enc_compute.jl	77.27% <0.00%> (-4.55%)	⬇️
lib/core/compile-opts.jl	86.84% <0.00%> (+2.63%)	⬆️

---

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update a1115cc...e234ecd. Read the comment docs.

[maleadt]

It looks like the gpuAddress property is available on older versions of Metal as well, despite being a private property, so we might just be able to support Monterey as well. If not, there's another possibility with the following hack:

const _kernel_cache = Dict{MtlDevice,MtlFunction}()
function gpuAddress(buf::MtlBuffer)
    # on Metal 2, we don't have a way to query the GPU address of a buffer. so instead,
    # we encode the buffer as the argument of a dummy kernel, and extract the encoded
    # address by peeking into the argument encoder's buffer. (note that this requires
    # an argument encoder, and thus a struct-valued argument, as the command encoder
    # doesn't have a buffer we can peek into).
    #
    # TODO: with Metal 3, the GPU address is a property of the buffer
    dev = current_device()
    fun = get!(_kernel_cache, dev) do
        code = """
            struct Struct {
                device void* buffer;
            };
            kernel void entry(device Struct*) { }"""
        lib = MtlLibrary(dev, code)
        MtlFunction(lib, "entry")
    end

    argenc = MtlArgumentEncoder(fun, 1)
    argbuf = alloc(Cchar, dev, sizeof(argenc), storage=Shared)
    MTL.assign_argument_buffer!(argenc, argbuf)
    set_buffer!(argenc, buf, 0, 0)
    Base.unsafe_load(Base.convert(Ptr{Ptr{Nothing}}, content(argbuf)), 1)
end

Putting this out here so that I can consider it if it's ever needed.