From 1cc21f4d83904a2859af4a15a63fad5bca8be89f Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Mon, 28 Sep 2026 22:20:20 +0200 Subject: [PATCH 1/7] Generate IR with LLVM.jl's @llvmgenerated and generate_llvmcall Local memory, invariant loads, the constant tables for randn, printf and the additional kernel arguments all built their IR with create_function/call_function around a lot of context and builder boilerplate. LLVM.jl 9.14's @llvmgenerated and generate_llvmcall derive the LLVM signature from the Julia one and verify the generated IR, so use those instead. printf needs the Julia types of its varargs, so it stays a @generated function that uses generate_llvmcall. --- Project.toml | 2 +- lib/intrinsics/Project.toml | 2 +- lib/intrinsics/src/memory.jl | 68 ++++++++++--------------- lib/intrinsics/src/printf.jl | 96 +++++++++++++++--------------------- src/device/array.jl | 48 ++++++------------ src/device/random.jl | 24 ++------- src/device/runtime.jl | 28 ++--------- 7 files changed, 93 insertions(+), 175 deletions(-) diff --git a/Project.toml b/Project.toml index cbaf7cdf..b4a6a2b3 100644 --- a/Project.toml +++ b/Project.toml @@ -35,7 +35,7 @@ GPUArrays = "11.2.1" GPUCompiler = "2.9" GPUToolbox = "3.3.2" KernelAbstractions = "0.9.38" -LLVM = "9.6" +LLVM = "9.14" LinearAlgebra = "1" OpenCL_jll = "=2024.10.24" Preferences = "1" diff --git a/lib/intrinsics/Project.toml b/lib/intrinsics/Project.toml index 8a6a2fbd..f179c9cb 100644 --- a/lib/intrinsics/Project.toml +++ b/lib/intrinsics/Project.toml @@ -18,7 +18,7 @@ SPIRVIntrinsicsSIMDExt = "SIMD" [compat] ExprTools = "0.1" GPUToolbox = "0.2, 0.3, 1, 2, 3" -LLVM = "9.1" +LLVM = "9.14" SIMD = "3.6" SpecialFunctions = "1.3, 2" julia = "1.10" diff --git a/lib/intrinsics/src/memory.jl b/lib/intrinsics/src/memory.jl index 9a17b447..487a7259 100644 --- a/lib/intrinsics/src/memory.jl +++ b/lib/intrinsics/src/memory.jl @@ -1,47 +1,31 @@ # local memory # get a pointer to local memory, with known (static) or zero length (dynamic) -@generated function emit_localmemory(::Type{T}, ::Val{len}=Val(0)) where {T,len} - Context() do ctx - # XXX: as long as LLVMPtr is emitted as i8*, it doesn't make sense to type the GV - eltyp = convert(LLVMType, LLVM.Int8Type()) - T_ptr = convert(LLVMType, LLVMPtr{T,AS.Workgroup}) - - # create a function - llvm_f, _ = create_function(T_ptr) - - # determine the array size: an array of a bits union stores a selector byte per - # element after the values - sz = len * sizeof(T) - Base.isbitsunion(T) && (sz += len) - - # create the global variable - mod = LLVM.parent(llvm_f) - gv_typ = LLVM.ArrayType(eltyp, sz) - gv = GlobalVariable(mod, gv_typ, "local_memory", AS.Workgroup) - if len > 0 - linkage!(gv, LLVM.API.LLVMInternalLinkage) - initializer!(gv, null(gv_typ)) - end - # TODO: Make the alignment configurable - align = 1 - for typ in Base.uniontypes(T) - typ.layout != C_NULL && (align = max(align, Base.datatype_alignment(typ))) - end - alignment!(gv, align) - - # generate IR - IRBuilder() do builder - entry = BasicBlock(llvm_f, "entry") - position!(builder, entry) - - ptr = gep!(builder, gv_typ, gv, [ConstantInt(0), ConstantInt(0)]) - - untyped_ptr = bitcast!(builder, ptr, T_ptr) - - ret!(builder, untyped_ptr) - end - - call_function(llvm_f, LLVMPtr{T,AS.Workgroup}) +@llvmgenerated builder function emit_localmemory(::Type{T}, + ::Val{len}=Val(0))::LLVMPtr{T,AS.Workgroup} where {T,len} + # XXX: as long as LLVMPtr is emitted as i8*, it doesn't make sense to type the GV + eltyp = LLVM.Int8Type() + T_ptr = convert(LLVMType, LLVMPtr{T,AS.Workgroup}) + + # determine the array size: an array of a bits union stores a selector byte per + # element after the values + sz = len * sizeof(T) + Base.isbitsunion(T) && (sz += len) + + # create the global variable + gv_typ = LLVM.ArrayType(eltyp, sz) + gv = GlobalVariable(current_module(builder), gv_typ, "local_memory", AS.Workgroup) + if len > 0 + linkage!(gv, LLVM.API.LLVMInternalLinkage) + initializer!(gv, null(gv_typ)) end + # TODO: Make the alignment configurable + align = 1 + for typ in Base.uniontypes(T) + typ.layout != C_NULL && (align = max(align, Base.datatype_alignment(typ))) + end + alignment!(gv, align) + + ptr = gep!(builder, gv_typ, gv, [ConstantInt(0), ConstantInt(0)]) + bitcast!(builder, ptr, T_ptr) end diff --git a/lib/intrinsics/src/printf.jl b/lib/intrinsics/src/printf.jl index 203a966d..3177c949 100644 --- a/lib/intrinsics/src/printf.jl +++ b/lib/intrinsics/src/printf.jl @@ -28,70 +28,54 @@ end arg_exprs = [:( argspec[$i] ) for i in 1:length(argspec)] arg_types = [argspec...] - Context() do ctx - T_void = LLVM.VoidType() + generate_llvmcall(Int32, Tuple{arg_types...}, arg_exprs...) do builder, args... T_int32 = LLVM.Int32Type() T_pint8 = LLVM.PointerType(LLVM.Int8Type(), AS.UniformConstant) - # create functions - param_types = LLVMType[convert(LLVMType, typ) for typ in arg_types] - llvm_f, _ = create_function(T_int32, param_types) - mod = LLVM.parent(llvm_f) - - IRBuilder() do builder - entry = BasicBlock(llvm_f, "entry") - position!(builder, entry) - - # `printf` needs to be invoked very specifically, e.g., the format string needs - # to be a pointer to a string, and arguments need to match exactly what is - # expected by the format string, so we cannot rely on how the arguments to this - # function have been passed in (by `llvmcall`). - T_actual_args = LLVMType[] - actual_args = LLVM.Value[] - for (_, (arg, argtyp)) in enumerate(zip(parameters(llvm_f), arg_types)) - if argtyp <: LLVMPtr - # passed as i8* - T,AS = argtyp.parameters - actual_typ = LLVM.PointerType(convert(LLVMType, T), AS) - actual_arg = bitcast!(builder, arg, actual_typ) - elseif argtyp <: Ptr - T = eltype(argtyp) - if T === Nothing - T = Int8 - end - actual_typ = LLVM.PointerType(convert(LLVMType, T)) - actual_arg = if value_type(arg) isa LLVM.PointerType - # passed as i8* or ptr - bitcast!(builder, arg, actual_typ) - else - # passed as i64 - inttoptr!(builder, arg, actual_typ) - end - elseif argtyp <: Bool - # passed as i8 - T = eltype(argtyp) - actual_typ = LLVM.Int1Type() - actual_arg = trunc!(builder, arg, actual_typ) + # `printf` needs to be invoked very specifically, e.g., the format string needs + # to be a pointer to a string, and arguments need to match exactly what is + # expected by the format string, so we cannot rely on how the arguments to this + # function have been passed in (by `llvmcall`). + T_actual_args = LLVMType[] + actual_args = LLVM.Value[] + for (arg, argtyp) in zip(args, arg_types) + if argtyp <: LLVMPtr + # passed as i8* + T,AS = argtyp.parameters + actual_typ = LLVM.PointerType(convert(LLVMType, T), AS) + actual_arg = bitcast!(builder, arg, actual_typ) + elseif argtyp <: Ptr + T = eltype(argtyp) + if T === Nothing + T = Int8 + end + actual_typ = LLVM.PointerType(convert(LLVMType, T)) + actual_arg = if value_type(arg) isa LLVM.PointerType + # passed as i8* or ptr + bitcast!(builder, arg, actual_typ) else - actual_typ = convert(LLVMType, argtyp) - actual_arg = arg + # passed as i64 + inttoptr!(builder, arg, actual_typ) end - push!(T_actual_args, actual_typ) - push!(actual_args, actual_arg) + elseif argtyp <: Bool + # passed as i8 + actual_typ = LLVM.Int1Type() + actual_arg = trunc!(builder, arg, actual_typ) + else + actual_typ = convert(LLVMType, argtyp) + actual_arg = arg end - - str = globalstring_ptr!(builder, String(fmt); addrspace=AS.UniformConstant) - - # invoke printf and return - printf_typ = LLVM.FunctionType(T_int32, [T_pint8]; vararg=true) - printf = LLVM.Function(mod, "printf", printf_typ) - push!(function_attributes(printf), EnumAttribute("nobuiltin")) - chars = call!(builder, printf_typ, printf, [str, actual_args...]) - - ret!(builder, chars) + push!(T_actual_args, actual_typ) + push!(actual_args, actual_arg) end - call_function(llvm_f, Int32, Tuple{arg_types...}, arg_exprs...) + str = globalstring_ptr!(builder, String(fmt); addrspace=AS.UniformConstant) + + # invoke printf and return + printf_typ = LLVM.FunctionType(T_int32, [T_pint8]; vararg=true) + printf = LLVM.Function(current_module(builder), "printf", printf_typ) + push!(function_attributes(printf), EnumAttribute("nobuiltin")) + call!(builder, printf_typ, printf, [str, actual_args...]) end end diff --git a/src/device/array.jl b/src/device/array.jl index 62466015..c961dc66 100644 --- a/src/device/array.jl +++ b/src/device/array.jl @@ -165,40 +165,24 @@ end # there is no SPIR-V equivalent of NVPTX's `ld.global.nc`, so instead of a dedicated # instruction we mark the load `!invariant.load`, which lets LLVM hoist it out of loops # and reorder it across stores to other objects. -@inline @generated function unsafe_invariant_load(ptr::LLVMPtr{T,AS}, i::I, - ::Val{align}) where {T,AS,I,align} +@inline function unsafe_invariant_load(ptr::LLVMPtr{T}, i::I, ::Val{align}) where {T,I,align} sizeof(T) == 0 && return T.instance - ispow2(align) || - return :(error("unsafe_invariant_load: alignment must be a power of 2, got $($align)")) - @dispose ctx=Context() begin - eltyp = convert(LLVMType, T) - T_idx = convert(LLVMType, I) - T_ptr = convert(LLVMType, ptr) - T_typed_ptr = LLVM.PointerType(eltyp, AS) - - llvm_f, _ = create_function(eltyp, LLVMType[T_ptr, T_idx]) - - @dispose builder=IRBuilder() begin - entry = BasicBlock(llvm_f, "entry") - position!(builder, entry) - base = if supports_typed_pointers(ctx) - bitcast!(builder, parameters(llvm_f)[1], T_typed_ptr) - else - parameters(llvm_f)[1] - end - gep = inbounds_gep!(builder, eltyp, base, [parameters(llvm_f)[2]]) - ld = load!(builder, eltyp, gep) - if AS != 0 - metadata(ld)[LLVM.MD_tbaa] = tbaa_addrspace(AS) - end - metadata(ld)[LLVM.MD_invariant_load] = MDNode(LLVM.Metadata[]) - alignment!(ld, align) - - ret!(builder, ld) - end - - call_function(llvm_f, T, Tuple{LLVMPtr{T,AS},I}, :ptr, :(i-one(I))) + ispow2(align) || error("unsafe_invariant_load: alignment must be a power of 2, got ", align) + return _unsafe_invariant_load(ptr, i - one(I), Val(align)) +end +@llvmgenerated builder function _unsafe_invariant_load(ptr::LLVMPtr{T,AS}, i::Integer, + ::Val{align})::T where {T,AS,align} + eltyp = convert(LLVMType, T) + if supports_typed_pointers(LLVM.context()) + ptr = bitcast!(builder, ptr, LLVM.PointerType(eltyp, AS)) + end + ld = load!(builder, eltyp, inbounds_gep!(builder, eltyp, ptr, [i])) + if AS != 0 + metadata(ld)[LLVM.MD_tbaa] = tbaa_addrspace(AS) end + metadata(ld)[LLVM.MD_invariant_load] = MDNode(LLVM.Metadata[]) + alignment!(ld, align) + ld end @device_function @inline function const_arrayref(A::CLDeviceArray{T}, index::Integer) where {T} diff --git a/src/device/random.jl b/src/device/random.jl index 88b7b225..53995242 100644 --- a/src/device/random.jl +++ b/src/device/random.jl @@ -150,36 +150,22 @@ end # a hacky method of exposing constant tables as constant GPU memory function emit_constant_array(name::Symbol, data::AbstractArray{T}) where {T} - @dispose ctx=Context() begin + generate_llvmcall(LLVMPtr{T,AS.UniformConstant}, Tuple{}) do builder T_val = convert(LLVMType, T) T_ptr = convert(LLVMType, LLVMPtr{T,AS.UniformConstant}) - # define function and get LLVM module - llvm_f, _ = create_function(T_ptr) - mod = LLVM.parent(llvm_f) - # create a global memory global variable # TODO: global_var alignment? T_global = LLVM.ArrayType(T_val, length(data)) # XXX: why can't we use a single name like emit_shmem - gv = GlobalVariable(mod, T_global, "gpu_$(name)_data", AS.UniformConstant) + gv = GlobalVariable(current_module(builder), T_global, "gpu_$(name)_data", + AS.UniformConstant) linkage!(gv, LLVM.API.LLVMInternalLinkage) initializer!(gv, ConstantArray(data)) alignment!(gv, 16) - # generate IR - @dispose builder=IRBuilder() begin - entry = BasicBlock(llvm_f, "entry") - position!(builder, entry) - - ptr = gep!(builder, T_global, gv, [ConstantInt(0), ConstantInt(0)]) - - untyped_ptr = bitcast!(builder, ptr, T_ptr) - - ret!(builder, untyped_ptr) - end - - call_function(llvm_f, LLVMPtr{T,AS.UniformConstant}) + ptr = gep!(builder, T_global, gv, [ConstantInt(0), ConstantInt(0)]) + bitcast!(builder, ptr, T_ptr) end end diff --git a/src/device/runtime.jl b/src/device/runtime.jl index c691b250..b997d31d 100644 --- a/src/device/runtime.jl +++ b/src/device/runtime.jl @@ -259,30 +259,10 @@ function additional_arg_intr(mod::LLVM.Module, T_state, name) end # run-time equivalent -function additional_arg_value(state, name) - @dispose ctx=Context() begin - T_state = convert(LLVMType, state) - - # create function - llvm_f, _ = create_function(T_state) - mod = LLVM.parent(llvm_f) - - # get intrinsic - state_intr = additional_arg_intr(mod, T_state, name) - state_intr_ft = function_type(state_intr) - - # generate IR - @dispose builder=IRBuilder() begin - entry = BasicBlock(llvm_f, "entry") - position!(builder, entry) - - val = call!(builder, state_intr_ft, state_intr, Value[], name) - - ret!(builder, val) - end - - call_function(llvm_f, state) - end +additional_arg_value(state, name) = generate_llvmcall(state, Tuple{}) do builder + T_state = convert(LLVMType, state) + state_intr = additional_arg_intr(current_module(builder), T_state, name) + call!(builder, function_type(state_intr), state_intr, Value[], name) end for name in [:random_keys, :random_counters] From 71935120b2ea8342965acf5c8f24f87bed2fcbcb Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 30 Sep 2026 17:33:05 +0200 Subject: [PATCH 2/7] Port to LLVM.jl 10 Import the IR and Build vocabularies next to LLVM.Interop, since `using LLVM` no longer exports the API, and use properties instead of the removed accessor functions (linkage, initializer, alignment, metadata, function attributes, module functions and globals, ...). Insertion points are explicit now: the block that initializes the RNG state is created with `LLVM.before` the entry block, and the builder positioned `LLVM.at_end` of it. That prologue gets its debug location from the builder only. The old `debuglocation!(builder, first(instructions(top_bb)))` copied the builder's location to the first instruction of the original entry block (despite what its docstring said), overwriting that instruction's location with the synthetic line 0, or clearing it when the kernel had no debug info, so it is dropped rather than ported. --- Project.toml | 2 +- lib/intrinsics/Project.toml | 2 +- lib/intrinsics/src/SPIRVIntrinsics.jl | 2 +- lib/intrinsics/src/memory.jl | 6 ++--- lib/intrinsics/src/printf.jl | 4 +-- src/OpenCL.jl | 2 +- src/compiler/compilation.jl | 35 +++++++++++++-------------- src/device/array.jl | 6 ++--- src/device/random.jl | 6 ++--- src/device/runtime.jl | 8 +++--- 10 files changed, 36 insertions(+), 37 deletions(-) diff --git a/Project.toml b/Project.toml index b4a6a2b3..3eab6c8b 100644 --- a/Project.toml +++ b/Project.toml @@ -35,7 +35,7 @@ GPUArrays = "11.2.1" GPUCompiler = "2.9" GPUToolbox = "3.3.2" KernelAbstractions = "0.9.38" -LLVM = "9.14" +LLVM = "10" LinearAlgebra = "1" OpenCL_jll = "=2024.10.24" Preferences = "1" diff --git a/lib/intrinsics/Project.toml b/lib/intrinsics/Project.toml index f179c9cb..6a0e4428 100644 --- a/lib/intrinsics/Project.toml +++ b/lib/intrinsics/Project.toml @@ -18,7 +18,7 @@ SPIRVIntrinsicsSIMDExt = "SIMD" [compat] ExprTools = "0.1" GPUToolbox = "0.2, 0.3, 1, 2, 3" -LLVM = "9.14" +LLVM = "10" SIMD = "3.6" SpecialFunctions = "1.3, 2" julia = "1.10" diff --git a/lib/intrinsics/src/SPIRVIntrinsics.jl b/lib/intrinsics/src/SPIRVIntrinsics.jl index ee432ca5..cf694e3e 100644 --- a/lib/intrinsics/src/SPIRVIntrinsics.jl +++ b/lib/intrinsics/src/SPIRVIntrinsics.jl @@ -1,6 +1,6 @@ module SPIRVIntrinsics -using LLVM, LLVM.Interop +using LLVM, LLVM.IR, LLVM.Build, LLVM.Interop using Core: LLVMPtr import ExprTools diff --git a/lib/intrinsics/src/memory.jl b/lib/intrinsics/src/memory.jl index 487a7259..7e796eb9 100644 --- a/lib/intrinsics/src/memory.jl +++ b/lib/intrinsics/src/memory.jl @@ -16,15 +16,15 @@ gv_typ = LLVM.ArrayType(eltyp, sz) gv = GlobalVariable(current_module(builder), gv_typ, "local_memory", AS.Workgroup) if len > 0 - linkage!(gv, LLVM.API.LLVMInternalLinkage) - initializer!(gv, null(gv_typ)) + gv.linkage = LLVM.Linkage.Internal + gv.initializer = null(gv_typ) end # TODO: Make the alignment configurable align = 1 for typ in Base.uniontypes(T) typ.layout != C_NULL && (align = max(align, Base.datatype_alignment(typ))) end - alignment!(gv, align) + gv.alignment = align ptr = gep!(builder, gv_typ, gv, [ConstantInt(0), ConstantInt(0)]) bitcast!(builder, ptr, T_ptr) diff --git a/lib/intrinsics/src/printf.jl b/lib/intrinsics/src/printf.jl index 3177c949..ae419f0b 100644 --- a/lib/intrinsics/src/printf.jl +++ b/lib/intrinsics/src/printf.jl @@ -50,7 +50,7 @@ end T = Int8 end actual_typ = LLVM.PointerType(convert(LLVMType, T)) - actual_arg = if value_type(arg) isa LLVM.PointerType + actual_arg = if arg.value_type isa LLVM.PointerType # passed as i8* or ptr bitcast!(builder, arg, actual_typ) else @@ -74,7 +74,7 @@ end # invoke printf and return printf_typ = LLVM.FunctionType(T_int32, [T_pint8]; vararg=true) printf = LLVM.Function(current_module(builder), "printf", printf_typ) - push!(function_attributes(printf), EnumAttribute("nobuiltin")) + push!(printf.function_attributes, EnumAttribute(:nobuiltin)) call!(builder, printf_typ, printf, [str, actual_args...]) end end diff --git a/src/OpenCL.jl b/src/OpenCL.jl index 8756602f..3e3ef109 100644 --- a/src/OpenCL.jl +++ b/src/OpenCL.jl @@ -2,7 +2,7 @@ module OpenCL using GPUCompiler import GPUToolbox -using LLVM, LLVM.Interop +using LLVM, LLVM.IR, LLVM.Build, LLVM.Interop using SPIRV_LLVM_Backend_jll, SPIRV_Tools_jll, spirv2clc_jll using Adapt using Reexport diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl index 413fd787..eadaa8b3 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -73,22 +73,22 @@ function GPUCompiler.finish_module!(@nospecialize(job::OpenCLCompilerJob), sg_size = job.config.params.sub_group_size if sg_size !== nothing - metadata(entry)["intel_reqd_sub_group_size"] = MDNode([ConstantInt(Int32(sg_size))]) + entry.metadata["intel_reqd_sub_group_size"] = MDNode([ConstantInt(Int32(sg_size))]) end # materialize the feature bitset for `has_feature`, if the kernel referenced it. A constant # initializer plus private linkage lets the optimizer fold the loads and drop the global, so it # never reaches SPIR-V. - if haskey(globals(mod), "__opencl_feature_bitset") - gv = globals(mod)["__opencl_feature_bitset"] - initializer!(gv, ConstantInt(LLVM.Int64Type(), job.config.params.features)) - linkage!(gv, LLVM.API.LLVMPrivateLinkage) - constant!(gv, true) + if haskey(mod.globals, "__opencl_feature_bitset") + gv = mod.globals["__opencl_feature_bitset"] + gv.initializer = ConstantInt(LLVM.Int64Type(), job.config.params.features) + gv.linkage = LLVM.Linkage.Private + gv.constant = true end # if this kernel uses our RNG, we should prime the shared state. # XXX: these transformations should really happen at the Julia IR level... - if haskey(functions(mod), "julia.opencl.random_keys") && job.config.kernel + if haskey(mod.functions, "julia.opencl.random_keys") && job.config.kernel # insert call to `initialize_rng_state` f = initialize_rng_state ft = typeof(f) @@ -102,16 +102,15 @@ function GPUCompiler.finish_module!(@nospecialize(job::OpenCLCompilerJob), GPUCompiler.deferred_codegen_jobs[id] = job # generate IR for calls to `deferred_codegen` and the resulting function pointer - top_bb = first(blocks(entry)) - bb = BasicBlock(top_bb, "initialize_rng") + top_bb = entry.entry + bb = BasicBlock(LLVM.before(top_bb), "initialize_rng") @dispose builder=IRBuilder() begin - position!(builder, bb) - subprogram = LLVM.subprogram(entry) + position!(builder, LLVM.at_end(bb)) + subprogram = entry.subprogram if subprogram !== nothing loc = DILocation(0, 0, subprogram) - debuglocation!(builder, loc) + builder.debug_location = loc end - debuglocation!(builder, first(instructions(top_bb))) # call the `deferred_codegen` marker function T_ptr = if LLVM.version() >= v"17" @@ -123,8 +122,8 @@ function GPUCompiler.finish_module!(@nospecialize(job::OpenCLCompilerJob), end T_id = convert(LLVMType, Int) deferred_codegen_ft = LLVM.FunctionType(T_ptr, [T_id]) - deferred_codegen = if haskey(functions(mod), "deferred_codegen") - functions(mod)["deferred_codegen"] + deferred_codegen = if haskey(mod.functions, "deferred_codegen") + mod.functions["deferred_codegen"] else LLVM.Function(mod, "deferred_codegen", deferred_codegen_ft) end @@ -139,7 +138,7 @@ function GPUCompiler.finish_module!(@nospecialize(job::OpenCLCompilerJob), br!(builder, top_bb) # note the use of the device-side RNG in this kernel - push!(function_attributes(entry), StringAttribute("julia.opencl.rng", "")) + push!(entry.function_attributes, StringAttribute("julia.opencl.rng", "")) end # XXX: put some of the above behind GPUCompiler abstractions @@ -247,8 +246,8 @@ function compile_to_obj(@nospecialize(job::CompilerJob)) JuliaContext() do ctx obj, meta = GPUCompiler.compile(:obj, job) - entry = LLVM.name(meta.entry) - device_rng = StringAttribute("julia.opencl.rng", "") in collect(function_attributes(meta.entry)) + entry = meta.entry.name + device_rng = StringAttribute("julia.opencl.rng", "") in collect(meta.entry.function_attributes) (; obj, entry, device_rng) end end diff --git a/src/device/array.jl b/src/device/array.jl index c961dc66..0f5f3997 100644 --- a/src/device/array.jl +++ b/src/device/array.jl @@ -178,10 +178,10 @@ end end ld = load!(builder, eltyp, inbounds_gep!(builder, eltyp, ptr, [i])) if AS != 0 - metadata(ld)[LLVM.MD_tbaa] = tbaa_addrspace(AS) + ld.metadata[LLVM.MD_tbaa] = tbaa_addrspace(AS) end - metadata(ld)[LLVM.MD_invariant_load] = MDNode(LLVM.Metadata[]) - alignment!(ld, align) + ld.metadata[LLVM.MD_invariant_load] = MDNode(LLVM.Metadata[]) + ld.alignment = align ld end diff --git a/src/device/random.jl b/src/device/random.jl index 53995242..a509ddcd 100644 --- a/src/device/random.jl +++ b/src/device/random.jl @@ -160,9 +160,9 @@ function emit_constant_array(name::Symbol, data::AbstractArray{T}) where {T} # XXX: why can't we use a single name like emit_shmem gv = GlobalVariable(current_module(builder), T_global, "gpu_$(name)_data", AS.UniformConstant) - linkage!(gv, LLVM.API.LLVMInternalLinkage) - initializer!(gv, ConstantArray(data)) - alignment!(gv, 16) + gv.linkage = LLVM.Linkage.Internal + gv.initializer = ConstantArray(data) + gv.alignment = 16 ptr = gep!(builder, T_global, gv, [ConstantInt(0), ConstantInt(0)]) bitcast!(builder, ptr, T_ptr) diff --git a/src/device/runtime.jl b/src/device/runtime.jl index b997d31d..162c3672 100644 --- a/src/device/runtime.jl +++ b/src/device/runtime.jl @@ -248,12 +248,12 @@ end # then get propagated across function calls to the caller. function additional_arg_intr(mod::LLVM.Module, T_state, name) - state_intr = if haskey(functions(mod), "julia.opencl.$name") - functions(mod)["julia.opencl.$name"] + state_intr = if haskey(mod.functions, "julia.opencl.$name") + mod.functions["julia.opencl.$name"] else LLVM.Function(mod, "julia.opencl.$name", LLVM.FunctionType(T_state)) end - push!(function_attributes(state_intr), EnumAttribute("readnone", 0)) + push!(state_intr.function_attributes, EnumAttribute(:readnone)) return state_intr end @@ -262,7 +262,7 @@ end additional_arg_value(state, name) = generate_llvmcall(state, Tuple{}) do builder T_state = convert(LLVMType, state) state_intr = additional_arg_intr(current_module(builder), T_state, name) - call!(builder, function_type(state_intr), state_intr, Value[], name) + call!(builder, state_intr.function_type, state_intr, Value[], name) end for name in [:random_keys, :random_counters] From 9d58b66bc0c461cd5b04defd57c8eb87dae5c3b2 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 30 Sep 2026 17:36:38 +0200 Subject: [PATCH 3/7] Use LLVM.jl's lookups instead of hand-written get-or-create code Look up or declare functions with `get!` on the module's function view, look up the feature bitset global with `get`, and check for the RNG attribute by its kind instead of collecting and comparing attributes. --- src/compiler/compilation.jl | 10 ++++------ src/device/runtime.jl | 11 ++++------- 2 files changed, 8 insertions(+), 13 deletions(-) diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl index eadaa8b3..aa47ced4 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -79,8 +79,8 @@ function GPUCompiler.finish_module!(@nospecialize(job::OpenCLCompilerJob), # materialize the feature bitset for `has_feature`, if the kernel referenced it. A constant # initializer plus private linkage lets the optimizer fold the loads and drop the global, so it # never reaches SPIR-V. - if haskey(mod.globals, "__opencl_feature_bitset") - gv = mod.globals["__opencl_feature_bitset"] + gv = get(mod.globals, "__opencl_feature_bitset", nothing) + if gv !== nothing gv.initializer = ConstantInt(LLVM.Int64Type(), job.config.params.features) gv.linkage = LLVM.Linkage.Private gv.constant = true @@ -122,9 +122,7 @@ function GPUCompiler.finish_module!(@nospecialize(job::OpenCLCompilerJob), end T_id = convert(LLVMType, Int) deferred_codegen_ft = LLVM.FunctionType(T_ptr, [T_id]) - deferred_codegen = if haskey(mod.functions, "deferred_codegen") - mod.functions["deferred_codegen"] - else + deferred_codegen = get!(mod.functions, "deferred_codegen") do LLVM.Function(mod, "deferred_codegen", deferred_codegen_ft) end fptr = call!(builder, deferred_codegen_ft, deferred_codegen, [ConstantInt(id)]) @@ -247,7 +245,7 @@ function compile_to_obj(@nospecialize(job::CompilerJob)) JuliaContext() do ctx obj, meta = GPUCompiler.compile(:obj, job) entry = meta.entry.name - device_rng = StringAttribute("julia.opencl.rng", "") in collect(meta.entry.function_attributes) + device_rng = haskey(meta.entry.function_attributes, "julia.opencl.rng") (; obj, entry, device_rng) end end diff --git a/src/device/runtime.jl b/src/device/runtime.jl index 162c3672..61f1d4c7 100644 --- a/src/device/runtime.jl +++ b/src/device/runtime.jl @@ -248,14 +248,11 @@ end # then get propagated across function calls to the caller. function additional_arg_intr(mod::LLVM.Module, T_state, name) - state_intr = if haskey(mod.functions, "julia.opencl.$name") - mod.functions["julia.opencl.$name"] - else - LLVM.Function(mod, "julia.opencl.$name", LLVM.FunctionType(T_state)) + get!(mod.functions, "julia.opencl.$name") do + state_intr = LLVM.Function(mod, "julia.opencl.$name", LLVM.FunctionType(T_state)) + push!(state_intr.function_attributes, EnumAttribute(:readnone)) + state_intr end - push!(state_intr.function_attributes, EnumAttribute(:readnone)) - - return state_intr end # run-time equivalent From 57792ef555b4a7528f4e877c06220edce4d082e9 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 30 Sep 2026 17:38:14 +0200 Subject: [PATCH 4/7] Generate IR with @llvmgenerated and clean up the generators The work-item built-ins, the barriers and the feature bitset were fixed IR strings that declare a global or function with attributes and load from or call it, and the accessor for the additional kernel arguments was a hand-written @generated function. Build them with @llvmgenerated instead, and declare the built-ins and the additional argument intrinsics with memory effects instead of `readnone`. Also make the invariant-load generator use an unconditional typed bitcast (which folds away with opaque pointers), `load!`'s `align` keyword and the public metadata kinds, derive the return type of the `deferred_codegen` declaration from Julia's lowering of `Ptr{Cvoid}` instead of a version decision tree, and drop printf's unused list of argument types. The IR generators of printf and of the RNG's constant tables were closures capturing the format string and the element type, which generate_llvmcall compiles for every one of them. Make them a named @nospecialize function that takes the format string and argument types as statically-known arguments, and a callable object with abstractly typed fields, so that they are compiled once. The generated IR is unchanged. --- lib/intrinsics/src/printf.jl | 96 ++++++++++++++------------- lib/intrinsics/src/synchronization.jl | 44 ++++++------ lib/intrinsics/src/work_item.jl | 46 +++++++------ src/compiler/capabilities.jl | 15 ++--- src/compiler/compilation.jl | 9 +-- src/device/array.jl | 12 ++-- src/device/random.jl | 45 ++++++++----- src/device/runtime.jl | 12 ++-- 8 files changed, 140 insertions(+), 139 deletions(-) diff --git a/lib/intrinsics/src/printf.jl b/lib/intrinsics/src/printf.jl index ae419f0b..28aa8066 100644 --- a/lib/intrinsics/src/printf.jl +++ b/lib/intrinsics/src/printf.jl @@ -26,57 +26,61 @@ end @generated function emit_printf(::Val{fmt}, argspec...) where {fmt} arg_exprs = [:( argspec[$i] ) for i in 1:length(argspec)] - arg_types = [argspec...] - - generate_llvmcall(Int32, Tuple{arg_types...}, arg_exprs...) do builder, args... - T_int32 = LLVM.Int32Type() - T_pint8 = LLVM.PointerType(LLVM.Int8Type(), AS.UniformConstant) - - # `printf` needs to be invoked very specifically, e.g., the format string needs - # to be a pointer to a string, and arguments need to match exactly what is - # expected by the format string, so we cannot rely on how the arguments to this - # function have been passed in (by `llvmcall`). - T_actual_args = LLVMType[] - actual_args = LLVM.Value[] - for (arg, argtyp) in zip(args, arg_types) - if argtyp <: LLVMPtr - # passed as i8* - T,AS = argtyp.parameters - actual_typ = LLVM.PointerType(convert(LLVMType, T), AS) - actual_arg = bitcast!(builder, arg, actual_typ) - elseif argtyp <: Ptr - T = eltype(argtyp) - if T === Nothing - T = Int8 - end - actual_typ = LLVM.PointerType(convert(LLVMType, T)) - actual_arg = if arg.value_type isa LLVM.PointerType - # passed as i8* or ptr - bitcast!(builder, arg, actual_typ) - else - # passed as i64 - inttoptr!(builder, arg, actual_typ) - end - elseif argtyp <: Bool - # passed as i8 - actual_typ = LLVM.Int1Type() - actual_arg = trunc!(builder, arg, actual_typ) + arg_types = Tuple{argspec...} + + # pass the format string and the argument types as statically-known arguments, so that + # the IR generator is compiled once instead of for every format string + generate_llvmcall(printf_ir, Int32, Tuple{Val{fmt}, Type{arg_types}, argspec...}, + Val(fmt), arg_types, arg_exprs...) +end + +function printf_ir(builder, fmt::Val, arg_types::Type{<:Tuple}, args...) + @nospecialize + T_int32 = LLVM.Int32Type() + T_pint8 = LLVM.PointerType(LLVM.Int8Type(), AS.UniformConstant) + + # `printf` needs to be invoked very specifically, e.g., the format string needs + # to be a pointer to a string, and arguments need to match exactly what is + # expected by the format string, so we cannot rely on how the arguments to this + # function have been passed in (by `llvmcall`). + actual_args = LLVM.Value[] + for (arg, argtyp) in zip(args, fieldtypes(arg_types)) + if argtyp <: LLVMPtr + # passed as i8* + T,AS = argtyp.parameters + actual_typ = LLVM.PointerType(convert(LLVMType, T), AS) + actual_arg = bitcast!(builder, arg, actual_typ) + elseif argtyp <: Ptr + T = eltype(argtyp) + if T === Nothing + T = Int8 + end + actual_typ = LLVM.PointerType(convert(LLVMType, T)) + actual_arg = if arg.value_type isa LLVM.PointerType + # passed as i8* or ptr + bitcast!(builder, arg, actual_typ) else - actual_typ = convert(LLVMType, argtyp) - actual_arg = arg + # passed as i64 + inttoptr!(builder, arg, actual_typ) end - push!(T_actual_args, actual_typ) - push!(actual_args, actual_arg) + elseif argtyp <: Bool + # passed as i8 + actual_typ = LLVM.Int1Type() + actual_arg = trunc!(builder, arg, actual_typ) + else + actual_arg = arg end + push!(actual_args, actual_arg) + end - str = globalstring_ptr!(builder, String(fmt); addrspace=AS.UniformConstant) + fmt_str = String(typeof(fmt).parameters[1]::Symbol) + str = globalstring_ptr!(builder, fmt_str; addrspace=AS.UniformConstant) - # invoke printf and return - printf_typ = LLVM.FunctionType(T_int32, [T_pint8]; vararg=true) - printf = LLVM.Function(current_module(builder), "printf", printf_typ) - push!(printf.function_attributes, EnumAttribute(:nobuiltin)) - call!(builder, printf_typ, printf, [str, actual_args...]) - end + # invoke printf and return + printf_typ = LLVM.FunctionType(T_int32, [T_pint8]; vararg=true) + printf = LLVM.Function(current_module(builder), "printf", printf_typ) + push!(printf.function_attributes, EnumAttribute(:nobuiltin)) + call!(builder, printf_typ, printf, [str, actual_args...]) end diff --git a/lib/intrinsics/src/synchronization.jl b/lib/intrinsics/src/synchronization.jl index e8cf6aef..c21af60b 100644 --- a/lib/intrinsics/src/synchronization.jl +++ b/lib/intrinsics/src/synchronization.jl @@ -34,38 +34,34 @@ end # XXX: is this even needed? Doesn't LLVM reconstruct these? # using the `@builtin_ccall` version causes validation issues. +# call a builtin that is `convergent`, i.e., that must not be made control-dependent on +# additional values +function convergent_call!(builder::IRBuilder, name::String, args::Vector{<:Value}) + ft = LLVM.FunctionType(LLVM.VoidType(), [arg.value_type for arg in args]) + f = LLVM.Function(current_module(builder), name, ft) + push!(f.function_attributes, EnumAttribute(:convergent)) + call!(builder, ft, f, args) + return +end + #@device_function @inline memory_barrier(scope, semantics) = # @builtin_ccall("__spirv_MemoryBarrier", Cvoid, (UInt32, UInt32), scope, semantics) @device_function memory_barrier(scope, semantics) = - Base.llvmcall((""" - declare void @_Z21__spirv_MemoryBarrierjj(i32, i32) #0 - define void @entry(i32 %scope, i32 %semantics) #1 { - call void @_Z21__spirv_MemoryBarrierjj(i32 %scope, i32 %semantics) - ret void - } - attributes #0 = { convergent } - attributes #1 = { alwaysinline } - """, "entry"), - Cvoid, Tuple{UInt32, UInt32}, convert(UInt32, scope), convert(UInt32, semantics)) + _memory_barrier(convert(UInt32, scope), convert(UInt32, semantics)) +@llvmgenerated builder function _memory_barrier(scope::UInt32, semantics::UInt32)::Nothing + convergent_call!(builder, "_Z21__spirv_MemoryBarrierjj", [scope, semantics]) +end #@device_function @inline control_barrier(execution_scope, memory_scope, memory_semantics) = # @builtin_ccall("__spirv_ControlBarrier", Cvoid, (UInt32, UInt32, UInt32), # execution_scope, memory_scope, memory_semantics) @device_function @inline control_barrier(execution_scope, memory_scope, memory_semantics) = - Base.llvmcall((""" - declare void @_Z22__spirv_ControlBarrierjjj(i32, i32, i32) #0 - define void @entry(i32 %execution, i32 %memory, i32 %semantics) #1 { - call void @_Z22__spirv_ControlBarrierjjj(i32 %execution, i32 %memory, i32 %semantics) - ret void - } - attributes #0 = { convergent } - attributes #1 = { alwaysinline } - """, "entry"), - Cvoid, - Tuple{UInt32, UInt32, UInt32}, - convert(UInt32, execution_scope), - convert(UInt32, memory_scope), - convert(UInt32, memory_semantics)) + _control_barrier(convert(UInt32, execution_scope), convert(UInt32, memory_scope), + convert(UInt32, memory_semantics)) +@llvmgenerated builder function _control_barrier(execution::UInt32, memory::UInt32, + semantics::UInt32)::Nothing + convergent_call!(builder, "_Z22__spirv_ControlBarrierjjj", [execution, memory, semantics]) +end ## OpenCL-compatible fence API diff --git a/lib/intrinsics/src/work_item.jl b/lib/intrinsics/src/work_item.jl index 5cf49bbf..9005ed0d 100644 --- a/lib/intrinsics/src/work_item.jl +++ b/lib/intrinsics/src/work_item.jl @@ -5,6 +5,25 @@ # NOTE: these functions now unsafely truncate to Int to avoid top bit checks. # we should probably use range metadata instead. +# load a built-in variable, which the SPIR-V back-end expects as an external global in the +# Input storage class +@llvmgenerated builder function builtin_variable(::Val{name}, ::Type{T})::T where {name,T} + T_val = convert(LLVMType, T) + gv = GlobalVariable(current_module(builder), T_val, String(name), AS.Input) + load!(builder, T_val, gv) +end + +# load a component of a built-in vector variable, by calling the function that the SPIR-V +# back-end lowers to a load and an extract of the component +@llvmgenerated builder function builtin_vector_variable(::Val{name}, idx::Int32)::UInt where {name} + ft = LLVM.FunctionType(convert(LLVMType, UInt), [idx.value_type]) + f = LLVM.Function(current_module(builder), String(name), ft) + push!(f.function_attributes, EnumAttribute(:nounwind)) + push!(f.function_attributes, EnumAttribute(:willreturn)) + f.memory_effects = MemoryEffects(:none) + call!(builder, ft, f, [idx]) +end + # 1D values for (julia_name, (spirv_name, julia_type, offset)) in [ # indices @@ -18,19 +37,11 @@ for (julia_name, (spirv_name, julia_type, offset)) in [ :get_max_sub_group_size => (:BuiltInSubgroupMaxSize, UInt32, 0), :get_num_sub_groups => (:BuiltInNumSubgroups, UInt32, 0), :get_enqueued_num_sub_groups => (:BuiltInNumEnqueuedSubgroups, UInt32, 0)] - gvar_name = Symbol("@__spirv_$(spirv_name)") - width = sizeof(julia_type) * 8 + gvar_name = Symbol("__spirv_$(spirv_name)") @eval begin export $julia_name @device_function $julia_name() = - Base.llvmcall( - $("""$gvar_name = external addrspace($(AS.Input)) global i$(width) - define i$(width) @entry() #0 { - %val = load i$(width), i$(width) addrspace($(AS.Input))* $gvar_name - ret i$(width) %val - } - attributes #0 = { alwaysinline } - """, "entry"), $julia_type, Tuple{}) % Int + $offset + builtin_variable(Val($(QuoteNode(gvar_name))), $julia_type) % Int + $offset end end @@ -52,20 +63,11 @@ for (julia_name, (spirv_name, offset)) in [ :get_enqueued_local_size => (:BuiltInEnqueuedWorkgroupSize, 0), :get_num_groups => (:BuiltInNumWorkgroups, 0)] fname = "__spirv_$(spirv_name)" - mangled = "_Z$(length(fname))$(fname)i" - push!(known_intrinsics, mangled) - width = Int === Int64 ? 64 : 32 + mangled = Symbol("_Z$(length(fname))$(fname)i") + push!(known_intrinsics, String(mangled)) @eval begin export $julia_name @device_function $julia_name(dimindx::Integer=1u32) = - Base.llvmcall( - $("""declare i$(width) @$(mangled)(i32) #0 - define i$(width) @entry(i32 %idx) #1 { - %val = call i$(width) @$(mangled)(i32 %idx) - ret i$(width) %val - } - attributes #0 = { nounwind readnone willreturn } - attributes #1 = { alwaysinline } - """, "entry"), UInt, Tuple{Int32}, (dimindx - 1u32) % Int32) % Int + $offset + builtin_vector_variable(Val($(QuoteNode(mangled))), (dimindx - 1u32) % Int32) % Int + $offset end end diff --git a/src/compiler/capabilities.jl b/src/compiler/capabilities.jl index 52562bad..631dcfd6 100644 --- a/src/compiler/capabilities.jl +++ b/src/compiler/capabilities.jl @@ -156,15 +156,12 @@ feature_supported(dev::cl.Device, name::Symbol) = feature_supported(device_featu # Load the feature bitset that `finish_module!` materializes as a module-scope constant. Once the # constant is in place the load folds away, so `has_feature` branches resolve at compile time. The # global uses the UniformConstant (2) storage class to stay valid SPIR-V if it ever survives. -@device_function @inline function feature_bitset() - Base.llvmcall( - ("""@__opencl_feature_bitset = external addrspace(2) global i64 - define i64 @entry() #0 { - %v = load i64, i64 addrspace(2)* @__opencl_feature_bitset - ret i64 %v - } - attributes #0 = { alwaysinline } - """, "entry"), UInt64, Tuple{}) +@device_function @inline feature_bitset() = _feature_bitset() +@llvmgenerated builder function _feature_bitset()::UInt64 + T = LLVM.Int64Type() + gv = GlobalVariable(current_module(builder), T, "__opencl_feature_bitset", + AS.UniformConstant) + load!(builder, T, gv) end export has_feature diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl index aa47ced4..1a2df70c 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -113,13 +113,8 @@ function GPUCompiler.finish_module!(@nospecialize(job::OpenCLCompilerJob), end # call the `deferred_codegen` marker function - T_ptr = if LLVM.version() >= v"17" - LLVM.PointerType() - elseif VERSION >= v"1.12.0-DEV.225" - LLVM.PointerType(LLVM.Int8Type()) - else - LLVM.Int64Type() - end + # (declared like GPUCompiler's `ccall("extern deferred_codegen", llvmcall, Ptr{Cvoid}, ...)`) + T_ptr = convert(LLVMType, Ptr{Cvoid}) T_id = convert(LLVMType, Int) deferred_codegen_ft = LLVM.FunctionType(T_ptr, [T_id]) deferred_codegen = get!(mod.functions, "deferred_codegen") do diff --git a/src/device/array.jl b/src/device/array.jl index 0f5f3997..c670478c 100644 --- a/src/device/array.jl +++ b/src/device/array.jl @@ -173,15 +173,13 @@ end @llvmgenerated builder function _unsafe_invariant_load(ptr::LLVMPtr{T,AS}, i::Integer, ::Val{align})::T where {T,AS,align} eltyp = convert(LLVMType, T) - if supports_typed_pointers(LLVM.context()) - ptr = bitcast!(builder, ptr, LLVM.PointerType(eltyp, AS)) - end - ld = load!(builder, eltyp, inbounds_gep!(builder, eltyp, ptr, [i])) + # `LLVMPtr` is an `i8*` with typed pointers (with opaque pointers, this cast folds away) + ptr = bitcast!(builder, ptr, LLVM.PointerType(eltyp, AS)) + ld = load!(builder, eltyp, inbounds_gep!(builder, eltyp, ptr, [i]); align) if AS != 0 - ld.metadata[LLVM.MD_tbaa] = tbaa_addrspace(AS) + ld.metadata[MD_tbaa] = tbaa_addrspace(AS) end - ld.metadata[LLVM.MD_invariant_load] = MDNode(LLVM.Metadata[]) - ld.alignment = align + ld.metadata[MD_invariant_load] = MDNode(LLVM.Metadata[]) ld end diff --git a/src/device/random.jl b/src/device/random.jl index a509ddcd..3ce97e66 100644 --- a/src/device/random.jl +++ b/src/device/random.jl @@ -150,23 +150,34 @@ end # a hacky method of exposing constant tables as constant GPU memory function emit_constant_array(name::Symbol, data::AbstractArray{T}) where {T} - generate_llvmcall(LLVMPtr{T,AS.UniformConstant}, Tuple{}) do builder - T_val = convert(LLVMType, T) - T_ptr = convert(LLVMType, LLVMPtr{T,AS.UniformConstant}) - - # create a global memory global variable - # TODO: global_var alignment? - T_global = LLVM.ArrayType(T_val, length(data)) - # XXX: why can't we use a single name like emit_shmem - gv = GlobalVariable(current_module(builder), T_global, "gpu_$(name)_data", - AS.UniformConstant) - gv.linkage = LLVM.Linkage.Internal - gv.initializer = ConstantArray(data) - gv.alignment = 16 - - ptr = gep!(builder, T_global, gv, [ConstantInt(0), ConstantInt(0)]) - bitcast!(builder, ptr, T_ptr) - end + generate_llvmcall(ConstantArrayIR(name, data), LLVMPtr{T,AS.UniformConstant}, Tuple{}) +end + +# the IR generator of `emit_constant_array`: a callable object with abstractly-typed fields +# instead of a closure, so that it is compiled once instead of for every element type +struct ConstantArrayIR + name::Symbol + data::AbstractArray +end + +function (gen::ConstantArrayIR)(builder) + (; name, data) = gen + T = eltype(data) + T_val = convert(LLVMType, T) + T_ptr = convert(LLVMType, LLVMPtr{T,AS.UniformConstant}) + + # create a global memory global variable + # TODO: global_var alignment? + T_global = LLVM.ArrayType(T_val, length(data)) + # XXX: why can't we use a single name like emit_shmem + gv = GlobalVariable(current_module(builder), T_global, "gpu_$(name)_data", + AS.UniformConstant) + gv.linkage = LLVM.Linkage.Internal + gv.initializer = ConstantArray(data) + gv.alignment = 16 + + ptr = gep!(builder, T_global, gv, [ConstantInt(0), ConstantInt(0)]) + bitcast!(builder, ptr, T_ptr) end for var in [:ki, :wi, :fi, :ke, :we, :fe] diff --git a/src/device/runtime.jl b/src/device/runtime.jl index 61f1d4c7..dee4ec57 100644 --- a/src/device/runtime.jl +++ b/src/device/runtime.jl @@ -250,19 +250,17 @@ end function additional_arg_intr(mod::LLVM.Module, T_state, name) get!(mod.functions, "julia.opencl.$name") do state_intr = LLVM.Function(mod, "julia.opencl.$name", LLVM.FunctionType(T_state)) - push!(state_intr.function_attributes, EnumAttribute(:readnone)) + state_intr.memory_effects = MemoryEffects(:none) state_intr end end # run-time equivalent -additional_arg_value(state, name) = generate_llvmcall(state, Tuple{}) do builder - T_state = convert(LLVMType, state) - state_intr = additional_arg_intr(current_module(builder), T_state, name) - call!(builder, state_intr.function_type, state_intr, Value[], name) +@llvmgenerated builder function additional_arg_value(::Type{T}, ::Val{name})::T where {T,name} + state_intr = additional_arg_intr(current_module(builder), convert(LLVMType, T), name) + call!(builder, state_intr.function_type, state_intr, Value[], String(name)) end for name in [:random_keys, :random_counters] - @eval @inline @generated $name() = - additional_arg_value(LLVMPtr{UInt32, AS.Workgroup}, $(String(name))) + @eval @inline $name() = additional_arg_value(LLVMPtr{UInt32, AS.Workgroup}, Val($(QuoteNode(name)))) end From 635aac3b8a2bed6510ab3665b37081270351bf72 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 30 Sep 2026 22:00:55 +0200 Subject: [PATCH 5/7] Widen the index of invariant loads to Int before the GEP The index was decremented in its own type and passed to getelementptr as is, which sign-extends narrower indices, so e.g. a UInt32 index of 0x80000001 addressed a negative offset. Convert it to Int in Julia first, like LLVM.jl's unsafe_load does. --- src/device/array.jl | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/src/device/array.jl b/src/device/array.jl index c670478c..85a7851c 100644 --- a/src/device/array.jl +++ b/src/device/array.jl @@ -165,12 +165,15 @@ end # there is no SPIR-V equivalent of NVPTX's `ld.global.nc`, so instead of a dedicated # instruction we mark the load `!invariant.load`, which lets LLVM hoist it out of loops # and reorder it across stores to other objects. -@inline function unsafe_invariant_load(ptr::LLVMPtr{T}, i::I, ::Val{align}) where {T,I,align} +# +# like `unsafe_load`, the index is widened to `Int` in Julia, where its signedness is known, +# because `getelementptr` sign-extends narrower indices. +@inline function unsafe_invariant_load(ptr::LLVMPtr{T}, i::Integer, ::Val{align}) where {T,align} sizeof(T) == 0 && return T.instance ispow2(align) || error("unsafe_invariant_load: alignment must be a power of 2, got ", align) - return _unsafe_invariant_load(ptr, i - one(I), Val(align)) + return _unsafe_invariant_load(ptr, Int(i) - 1, Val(align)) end -@llvmgenerated builder function _unsafe_invariant_load(ptr::LLVMPtr{T,AS}, i::Integer, +@llvmgenerated builder function _unsafe_invariant_load(ptr::LLVMPtr{T,AS}, i::Int, ::Val{align})::T where {T,AS,align} eltyp = convert(LLVMType, T) # `LLVMPtr` is an `i8*` with typed pointers (with opaque pointers, this cast folds away) From ca4f753d12a5c663d1c0e95aff75a85ceb182bef Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Thu, 1 Oct 2026 14:13:32 +0200 Subject: [PATCH 6/7] Dispose of the compiled IR once it has been inspected GPUCompiler now hands the module that Julia's code generator produced to the caller of compile: on Julia 1.11 and later it moves the module out of the native code, and on 1.10 (LLVM 15, which can't do that) it returns the module for the caller to consume or dispose of. compile_to_obj only read the entry point's name and attributes, and never disposed of the module. That leaked it: up to Julia 1.13, the native-code descriptor is never freed and keeps its thread-safe module, and with it the context, alive, so the module isn't freed together with the context either. memcheck reported every compiled module as leaked. Dispose of the IR with @dispose inside the JuliaContext block, once the entry point has been inspected, so that it is also disposed of when the inspection throws. The test that compiles kernels for a device without Float64 discarded the IR returned by compile(:llvm) as well; dispose of it there too. --- src/compiler/compilation.jl | 10 +++++++--- test/intrinsics.jl | 7 ++++++- 2 files changed, 13 insertions(+), 4 deletions(-) diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl index 1a2df70c..4cbe085c 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -239,9 +239,13 @@ function compile_to_obj(@nospecialize(job::CompilerJob)) JuliaContext() do ctx obj, meta = GPUCompiler.compile(:obj, job) - entry = meta.entry.name - device_rng = haskey(meta.entry.function_attributes, "julia.opencl.rng") - (; obj, entry, device_rng) + + # we own the IR: inspect it, then dispose of it + @dispose ir=meta.ir begin + entry = meta.entry.name + device_rng = haskey(meta.entry.function_attributes, "julia.opencl.rng") + (; obj, entry, device_rng) + end end end diff --git a/test/intrinsics.jl b/test/intrinsics.jl index 590443de..4c1b41d8 100644 --- a/test/intrinsics.jl +++ b/test/intrinsics.jl @@ -391,6 +391,11 @@ end (inv, ComplexF32, (ComplexF32,))) tt = Tuple{CLDeviceArray{T,0,AS.CrossWorkgroup}, typeof(f), args...} job = GPUCompiler.CompilerJob(GPUCompiler.methodinstance(typeof(kernel), tt), config) - @test GPUCompiler.JuliaContext(_ -> GPUCompiler.compile(:llvm, job)) !== nothing + # the IR belongs to the caller of `compile` + @test GPUCompiler.JuliaContext() do _ + ir, meta = GPUCompiler.compile(:llvm, job) + GPUCompiler.LLVM.dispose(ir) + meta !== nothing + end end end From e54feca8aef1c62c423962475b08ca8b6090a342 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Thu, 1 Oct 2026 20:04:07 +0200 Subject: [PATCH 7/7] Bump SPIRVIntrinsics to 1.2.0 and require it SPIRVIntrinsics 1.2.0 requires LLVM.jl 10, so OpenCL.jl, which uses it through lib/intrinsics, requires this version too. SPIRVIntrinsics is released from this repository, and oneAPI and KernelAbstractions need the new version as well. --- Project.toml | 2 +- lib/intrinsics/Project.toml | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/Project.toml b/Project.toml index 3eab6c8b..ec86c38c 100644 --- a/Project.toml +++ b/Project.toml @@ -44,7 +44,7 @@ Random = "1" Random123 = "1.7.1" RandomNumbers = "1.6.0" Reexport = "1" -SPIRVIntrinsics = "1.1" +SPIRVIntrinsics = "1.2" SPIRV_LLVM_Backend_jll = "23" SPIRV_Tools_jll = "2025.1" StaticArrays = "1" diff --git a/lib/intrinsics/Project.toml b/lib/intrinsics/Project.toml index 6a0e4428..57f56bd7 100644 --- a/lib/intrinsics/Project.toml +++ b/lib/intrinsics/Project.toml @@ -1,7 +1,7 @@ name = "SPIRVIntrinsics" uuid = "71d1d633-e7e8-4a92-83a1-de8814b09ba8" authors = ["Tim Besard "] -version = "1.1.4" +version = "1.2.0" [deps] ExprTools = "e2ba6199-217a-4e67-a87a-7c52f15ade04"