Autorelease changes lead to use after free with errors

[maleadt]

The following MWE fails under validation on an M3:

using Metal

function kernel(arr_sin, arr_cos)
    idx = thread_position_in_grid_1d()
    arr_sin[idx], arr_cos[idx] = sincos(arr_cos[idx])
    return
end

a = MtlArray{Float32}(undef, 1)
b = MtlArray{Float32}(undef, 1)
@metal kernel(a, b)

Normallly this renders as:

❯ MTL_SHADER_VALIDATION=1 jl wip.jl
2024-03-04 15:20:29.314 julia[64445:10376600] Metal GPU Validation Enabled
ERROR: LoadError: Compilation to native code failed; see below for details.
If you think this is a bug, please file an issue and attach /var/folders/wc/z_cydqs118735x03cvsf96tr0000gn/T/jl_SzYngbwtO5.metallib
Stacktrace:
  [1] error(s::String)
    @ Base ./error.jl:35
  [2] link(job::GPUCompiler.CompilerJob, compiled::@NamedTuple{image::Vector{UInt8}, entry::String}; return_function::Bool)
    @ Metal ~/Julia/pkg/Metal/src/compiler/compilation.jl:123
  [3] link(job::GPUCompiler.CompilerJob, compiled::@NamedTuple{image::Vector{UInt8}, entry::String})
    @ Metal ~/Julia/pkg/Metal/src/compiler/compilation.jl:110
  [4] actual_compilation(cache::Dict{Any, Any}, src::Core.MethodInstance, world::UInt64, cfg::GPUCompiler.CompilerConfig{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, compiler::typeof(Metal.compile), linker::typeof(Metal.link))
    @ GPUCompiler ~/Julia/pkg/GPUCompiler/src/execution.jl:132
  [5] cached_compilation(cache::Dict{Any, Any}, src::Core.MethodInstance, cfg::GPUCompiler.CompilerConfig{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, compiler::Function, linker::Function)
    @ GPUCompiler ~/Julia/pkg/GPUCompiler/src/execution.jl:103
  [6] macro expansion
    @ ~/Julia/pkg/Metal/src/compiler/execution.jl:185 [inlined]
  [7] macro expansion
    @ ./lock.jl:267 [inlined]
  [8] mtlfunction(f::typeof(kernel), tt::Type{Tuple{MtlDeviceVector{Float32, 1}, MtlDeviceVector{Float32, 1}}}; name::Nothing, kwargs::@Kwargs{})
    @ Metal ~/Julia/pkg/Metal/src/compiler/execution.jl:180
  [9] mtlfunction(f::typeof(kernel), tt::Type{Tuple{MtlDeviceVector{Float32, 1}, MtlDeviceVector{Float32, 1}}})
    @ Metal ~/Julia/pkg/Metal/src/compiler/execution.jl:178
 [10] top-level scope
    @ ~/Julia/pkg/Metal/src/compiler/execution.jl:85
in expression starting at /Users/tim/Julia/pkg/Metal/wip.jl:11

caused by: NSError: Compiler encountered an internal error (AGXMetalG15X_B0, code 3)
Stacktrace:
  [1] MTLComputePipelineState(dev::Metal.MTL.MTLDeviceInstance, fun::Metal.MTL.MTLFunctionInstance)
    @ Metal.MTL ~/Julia/pkg/Metal/lib/mtl/compute_pipeline.jl:60
  [2] link(job::GPUCompiler.CompilerJob, compiled::@NamedTuple{image::Vector{UInt8}, entry::String}; return_function::Bool)
    @ Metal ~/Julia/pkg/Metal/src/compiler/compilation.jl:115
  [3] link(job::GPUCompiler.CompilerJob, compiled::@NamedTuple{image::Vector{UInt8}, entry::String})
    @ Metal ~/Julia/pkg/Metal/src/compiler/compilation.jl:110
  [4] actual_compilation(cache::Dict{Any, Any}, src::Core.MethodInstance, world::UInt64, cfg::GPUCompiler.CompilerConfig{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, compiler::typeof(Metal.compile), linker::typeof(Metal.link))
    @ GPUCompiler ~/Julia/pkg/GPUCompiler/src/execution.jl:132
  [5] cached_compilation(cache::Dict{Any, Any}, src::Core.MethodInstance, cfg::GPUCompiler.CompilerConfig{GPUCompiler.MetalCompilerTarget, Metal.MetalCompilerParams}, compiler::Function, linker::Function)
    @ GPUCompiler ~/Julia/pkg/GPUCompiler/src/execution.jl:103
  [6] macro expansion
    @ ~/Julia/pkg/Metal/src/compiler/execution.jl:185 [inlined]
  [7] macro expansion
    @ ./lock.jl:267 [inlined]
  [8] mtlfunction(f::typeof(kernel), tt::Type{Tuple{MtlDeviceVector{Float32, 1}, MtlDeviceVector{Float32, 1}}}; name::Nothing, kwargs::@Kwargs{})
    @ Metal ~/Julia/pkg/Metal/src/compiler/execution.jl:180
  [9] mtlfunction(f::typeof(kernel), tt::Type{Tuple{MtlDeviceVector{Float32, 1}, MtlDeviceVector{Float32, 1}}})
    @ Metal ~/Julia/pkg/Metal/src/compiler/execution.jl:178
 [10] top-level scope
    @ ~/Julia/pkg/Metal/src/compiler/execution.jl:85

After #294, the caused by leads to a segfault, presumable due to an UAF:

caused by:
[64465] signal (11.2): Segmentation fault: 11
in expression starting at none:0
objc_msgSend at /usr/lib/libobjc.A.dylib (unknown line)
getproperty at /Users/tim/Julia/pkg/ObjectiveC/src/syntax.jl:163 [inlined]
showerror at /Users/tim/Julia/pkg/ObjectiveC/src/foundation.jl:287
Allocations: 13979919 (Pool: 13963570; Big: 16349); GC: 19

[64465] signal (10.1): Bus error: 10
in expression starting at none:0
lookUpImpOrForward at /usr/lib/libobjc.A.dylib (unknown line)
Allocations: 13979919 (Pool: 13963570; Big: 16349); GC: 19