ERROR: old function still has uses (via a constant expr)

[alecloudenback]

In trying to work around JuliaGPU/Metal.jl#570 I've been running a bunch of code that ends up hitting that compile error in JuliaGPU/Metal.jl#570. Somewhere along the way it seems (based on the GPUCompiler src code) like I've messed up the LLVM cache somehow.

Maybe related to JuliaGPU/Metal.jl#570, but here no compiler error or multiplication.

MWE:

function w(out, model_object)
    i = thread_position_in_grid_1d()
    return if i <= 100
        for t in 1:20
            out[t, i] = t // model_object[i]
        end
    end
end

let
    c = Metal.fill(1, 100)

    out = Metal.zeros(100)
    threads = 16
    groups = cld.(size(c,1), threads)

    Metal.@sync @metal threads = 16 groups = groups w(out, c)
    out
end

If the line out[t, i] = t // model_object[I] is changed to either of the following, the code still works:

• out[t, i] = t
• out[t,i] = 1 // model_object[I]
• out[t, i] = t * model_object[i]
• out[t, i] = t // 12

So it seems like it's the combination of the loop variable t with getting an index from a vector.

Stacktrace:

ERROR: old function still has uses (via a constant expr)
Stacktrace:
  [1] error(s::String)
    @ Base ./error.jl:35
  [2] add_global_address_spaces!(job::GPUCompiler.CompilerJob, mod::LLVM.Module, entry::LLVM.Function)
    @ GPUCompiler ~/.julia/packages/GPUCompiler/OGnEB/src/metal.jl:413
  [3] finish_ir!(job::GPUCompiler.CompilerJob{GPUCompiler.MetalCompilerTarget}, mod::LLVM.Module, entry::LLVM.Function)
    @ GPUCompiler ~/.julia/packages/GPUCompiler/OGnEB/src/metal.jl:166
  [4] finish_ir!(job::GPUCompiler.CompilerJob{…}, mod::LLVM.Module, entry::LLVM.Function)
    @ Metal ~/.julia/packages/Metal/XF4Hj/src/compiler/compilation.jl:14
  [5] macro expansion
    @ ~/.julia/packages/GPUCompiler/OGnEB/src/driver.jl:283 [inlined]
  [6] 
    @ GPUCompiler ~/.julia/packages/GPUCompiler/OGnEB/src/utils.jl:110
  [7] 
    @ GPUCompiler ~/.julia/packages/GPUCompiler/OGnEB/src/driver.jl:100
  [8] codegen
    @ ~/.julia/packages/GPUCompiler/OGnEB/src/driver.jl:82 [inlined]
  [9] compile(target::Symbol, job::GPUCompiler.CompilerJob; kwargs::@Kwargs{})
    @ GPUCompiler ~/.julia/packages/GPUCompiler/OGnEB/src/driver.jl:79
 [10] compile
    @ ~/.julia/packages/GPUCompiler/OGnEB/src/driver.jl:74 [inlined]
 [11] (::Metal.var"#155#163"{GPUCompiler.CompilerJob{…}})(ctx::LLVM.Context)
    @ Metal ~/.julia/packages/Metal/XF4Hj/src/compiler/compilation.jl:108
 [12] JuliaContext(f::Metal.var"#155#163"{GPUCompiler.CompilerJob{…}}; kwargs::@Kwargs{})
    @ GPUCompiler ~/.julia/packages/GPUCompiler/OGnEB/src/driver.jl:34
 [13] JuliaContext(f::Function)
    @ GPUCompiler ~/.julia/packages/GPUCompiler/OGnEB/src/driver.jl:25
 [14] macro expansion
    @ ~/.julia/packages/Metal/XF4Hj/src/compiler/compilation.jl:107 [inlined]
 [15] macro expansion
    @ ~/.julia/packages/ObjectiveC/TgrW6/src/os.jl:264 [inlined]
 [16] compile(job::GPUCompiler.CompilerJob)
    @ Metal ~/.julia/packages/Metal/XF4Hj/src/compiler/compilation.jl:105
 [17] actual_compilation(cache::Dict{…}, src::Core.MethodInstance, world::UInt64, cfg::GPUCompiler.CompilerConfig{…}, compiler::typeof(Metal.compile), linker::typeof(Metal.link))
    @ GPUCompiler ~/.julia/packages/GPUCompiler/OGnEB/src/execution.jl:237
 [18] cached_compilation(cache::Dict{…}, src::Core.MethodInstance, cfg::GPUCompiler.CompilerConfig{…}, compiler::Function, linker::Function)
    @ GPUCompiler ~/.julia/packages/GPUCompiler/OGnEB/src/execution.jl:151
 [19] macro expansion
    @ ~/.julia/packages/Metal/XF4Hj/src/compiler/execution.jl:189 [inlined]
 [20] macro expansion
    @ ./lock.jl:273 [inlined]
 [21] mtlfunction(f::typeof(w), tt::Type{Tuple{MtlDeviceVector{…}, MtlDeviceVector{…}}}; name::Nothing, kwargs::@Kwargs{})
    @ Metal ~/.julia/packages/Metal/XF4Hj/src/compiler/execution.jl:184
 [22] mtlfunction(f::typeof(w), tt::Type{Tuple{MtlDeviceVector{Float32, 1}, MtlDeviceVector{Int64, 1}}})
    @ Metal ~/.julia/packages/Metal/XF4Hj/src/compiler/execution.jl:182
 [23] macro expansion
    @ ~/.julia/packages/Metal/XF4Hj/src/compiler/execution.jl:85 [inlined]
 [24] macro expansion
    @ ~/.julia/packages/Metal/XF4Hj/src/utilities.jl:10 [inlined]
 [25] top-level scope
    @ ~/prog/scratchpad/myscript/script.jl:118
Some type information was truncated. Use `show(err)` to see complete types.

julia> Metal.versioninfo()
macOS 15.3.2, Darwin 24.3.0

Toolchain:
- Julia: 1.11.4
- LLVM: 16.0.6

Julia packages: 
- Metal.jl: 1.5.1
- GPUArrays: 11.2.2
- GPUCompiler: 1.2.0
- KernelAbstractions: 0.9.34
- ObjectiveC: 3.4.1
- LLVM: 9.2.0
- LLVMDowngrader_jll: 0.6.0+0

1 device:
- Apple M4 Max (108.578 MiB allocated)

I tried the following, but still get the error:

• restarting Julia session
• deleting the jl_* files in /var/folders/hw/bycsc7f52zvfzlh61rydj85r0000gn/T/
• using Metal.jl main branch
• removing Metal, doing Pkg gc, and then re-adding
• deleting .Julia/v1.11/Metal/

[christiangnrd]

Are you able to provide a MWE?

[alecloudenback]

Updated to have a MWE

[alecloudenback]

Per the suggestion in JuliaGPU/Metal.jl#570, I tested

out[t, i] = Rational{Int32}(t, model_object[I]) and this still has the same error.

[christiangnrd]

This is probably related to JuliaGPU/Metal.jl#69

[maleadt]

#678 should fix the assertion; please verify.

Note that the kernel doesn't successfully compile, as it contains dynamic code:

ERROR: InvalidIRError: compiling MethodInstance for w(::MtlDeviceVector{Float32, 1}, ::MtlDeviceVector{Int64, 1}) resulted in invalid LLVM IR
Reason: unsupported dynamic function invocation (call to __throw_rational_argerror_zero(T) @ Base rational.jl:30)
Stacktrace:
 [1] Rational
   @ ./rational.jl:32
 [2] Rational
   @ ./rational.jl:48
 [3] //
   @ ./rational.jl:84
 [4] w
   @ ./REPL[3]:5