From 675036dd01da6f86e61d8a071a90df102f2f8a83 Mon Sep 17 00:00:00 2001 From: Ian Butterworth Date: Mon, 26 Oct 2026 21:17:19 -0400 Subject: [PATCH 1/2] pocl: Define the device runtime as overlays so sysimages load The POCL device runtime was defined as ordinary methods whose bodies declare device intrinsics such as `julia.gpu.state_getter`. Sysimage builds compile those methods for the host: Julia 1.12 and later compile every method with a compileable signature, and Julia 1.11 keeps the native code instances that inferring `initialize_rng_state` left behind. A sysimage is loaded with RTLD_NOW, so the undefined intrinsics stop it from loading. Define the device-only runtime functions with `@device_function`, so the host only sees a method that throws and the real body lives in the overlay table, as SPIRVIntrinsics does for its built-ins. `malloc` and `report_oom` use no intrinsics and stay host methods. Look up `initialize_rng_state` through the job's method tables and call it as a void function, instead of inferring its return type natively. Test that no method of the package reaches device code on the host, with GPUToolbox's `host_leaks`. This needs SPIRVIntrinsics 1.6, whose `@device_function` defines the body as a plain overlay, and GPUToolbox 3.5. Co-Authored-By: Claude Opus 5.5 --- Project.toml | 4 ++-- src/pocl/compiler/compilation.jl | 14 ++++++++------ src/pocl/device/random.jl | 12 +++++++----- src/pocl/device/runtime.jl | 24 +++++++++++++++--------- test/Project.toml | 1 + test/runtests.jl | 7 +++++++ 6 files changed, 40 insertions(+), 22 deletions(-) diff --git a/Project.toml b/Project.toml index e9443f4a9..5afd14bed 100644 --- a/Project.toml +++ b/Project.toml @@ -50,7 +50,7 @@ AMDGPU = "2" Adapt = "0.4, 1.0, 2.0, 3.0, 4" Atomix = "1.5" GPUCompiler = "2.18" -GPUToolbox = "3.4.1" +GPUToolbox = "3.5" KernelInterface = "0.4" LLVM = "10" LinearAlgebra = "1.6" @@ -61,7 +61,7 @@ Printf = "<0.0.1, 1" Random = "1" Random123 = "1.7.1" RandomNumbers = "1.6.0" -SPIRVIntrinsics = "1.4" +SPIRVIntrinsics = "1.6" SPIRV_LLVM_Backend_jll = "23" SPIRV_Tools_jll = "2024.4, 2025.1" SparseArrays = "<0.0.1, 1.6" diff --git a/src/pocl/compiler/compilation.jl b/src/pocl/compiler/compilation.jl index 5649b5ac1..4faa52c6b 100644 --- a/src/pocl/compiler/compilation.jl +++ b/src/pocl/compiler/compilation.jl @@ -79,8 +79,11 @@ function GPUCompiler.finish_module!( ft = typeof(f) tt = Tuple{} - # create a deferred compilation job for `initialize_rng_state` - src = methodinstance(ft, tt, job.world) + # create a deferred compilation job for `initialize_rng_state`. it is an overlay, so + # look it up in the job's method tables. + src = GPUCompiler.generic_methodinstance( + ft, tt, job.world; method_table_view = GPUCompiler.method_table_view(job) + ) cfg = CompilerConfig(job.config; kernel = false, name = nothing) job = CompilerJob(src, cfg, job.world) id = length(GPUCompiler.deferred_codegen_jobs) + 1 @@ -107,10 +110,9 @@ function GPUCompiler.finish_module!( end fptr = call!(builder, deferred_codegen_ft, deferred_codegen, [ConstantInt(id)]) - # call the `initialize_rng_state` function - rt = Core.Compiler.return_type(f, tt) - llvm_rt = convert(LLVMType, rt) - llvm_ft = LLVM.FunctionType(llvm_rt) + # call the `initialize_rng_state` function, which returns nothing. inferring its + # return type natively would leave host code for it in sysimages. + llvm_ft = LLVM.FunctionType(LLVM.VoidType()) fptr = inttoptr!(builder, fptr, LLVM.PointerType(llvm_ft)) call!(builder, llvm_ft, fptr) br!(builder, top_bb) diff --git a/src/pocl/device/random.jl b/src/pocl/device/random.jl index b3f8f52f0..f7ebeaec3 100644 --- a/src/pocl/device/random.jl +++ b/src/pocl/device/random.jl @@ -4,24 +4,26 @@ using Random import RandomNumbers # local memory with the actual seed, per subgroup, set by `initialize_rng_state`` or overridden by calling `seed!` -@inline function global_random_keys() +@device_function @inline function global_random_keys() n = get_num_sub_groups() ptr = random_keys()::LLVMPtr{UInt32, AS.Workgroup} return CLDeviceArray{UInt32, 1, AS.Workgroup}((n,), ptr) end # local memory with per-subgroup counters, incremented when generating numbers -@inline function global_random_counters() +@device_function @inline function global_random_counters() n = get_num_sub_groups() ptr = random_counters()::LLVMPtr{UInt32, AS.Workgroup} return CLDeviceArray{UInt32, 1, AS.Workgroup}((n,), ptr) end -# initialization function, called automatically at the start of each kernel -function initialize_rng_state() +# initialization function, called automatically at the start of each kernel. +# it must return nothing, as the kernel calls it as a void function. +@device_function function initialize_rng_state() subgroup_id = get_sub_group_id() @inbounds global_random_keys()[subgroup_id] = kernel_state().random_seed - return @inbounds global_random_counters()[subgroup_id] = 0 + @inbounds global_random_counters()[subgroup_id] = 0 + return end # generators diff --git a/src/pocl/device/runtime.jl b/src/pocl/device/runtime.jl index 046cfb616..6a6bd5288 100644 --- a/src/pocl/device/runtime.jl +++ b/src/pocl/device/runtime.jl @@ -20,15 +20,19 @@ end return :(reinterpret(LLVMPtr{$T, AS.CrossWorkgroup}, kernel_state().exception_info + $offset)) end +# device-only functions are overlays, with a host method that throws. sysimage builds compile +# host methods, and a sysimage that references device intrinsics fails to load +# (JuliaGPU/GPUCompiler.jl#611). + # the record is shared by all work-items on the device, so these need device scope, which # UnsafeAtomics' primitives emit as given (SPIRVIntrinsics' atomics are relaxed) -@inline function atomic_store_device!(ptr::LLVMPtr{Int32, AS.CrossWorkgroup}, val::Int32) +@device_function @inline function atomic_store_device!(ptr::LLVMPtr{Int32, AS.CrossWorkgroup}, val::Int32) UnsafeAtomics.Internal.llvm_store!( ptr, val, Val(:release), Val(:device), Val(false), Val(sizeof(Int32)), Val(()) ) return end -@inline function atomic_cas_device!(ptr::LLVMPtr{Int32, AS.CrossWorkgroup}, cmp::Int32, val::Int32) +@device_function @inline function atomic_cas_device!(ptr::LLVMPtr{Int32, AS.CrossWorkgroup}, cmp::Int32, val::Int32) (; old) = UnsafeAtomics.Internal.llvm_cmpxchg!( ptr, cmp, val, Val(:acq_rel), Val(:acquire), Val(:device), Val(false), Val(false), Val(sizeof(Int32)), Val(()) @@ -40,7 +44,7 @@ end # would be interleaved. it can do so over several calls (e.g., a quirk printing the exception, # and GPUCompiler printing a backtrace). returns 1 when the output was just claimed, 2 when # it was claimed by an earlier call of this work-item, and 0 when it belongs to another. -@noinline function claim_output() +@device_function @noinline function claim_output() lock = exception_field(Val(:output_lock)) me = (get_global_id(1), get_global_id(2), get_global_id(3), 0) state = atomic_cas_device!(lock, Int32(0), Int32(1)) @@ -55,7 +59,7 @@ end return state == 2 && unsafe_load(exception_field(Val(:owner))) == me ? 2 : 0 end -function report_exception(ex) +@device_function function report_exception(ex) claim = claim_output() if claim == 1 SPIRVIntrinsics.@printf( @@ -68,7 +72,7 @@ function report_exception(ex) return end -function report_exception_name(ex) +@device_function function report_exception_name(ex) claim = claim_output() if claim == 1 SPIRVIntrinsics.@printf( @@ -81,7 +85,7 @@ function report_exception_name(ex) return end -function report_exception_frame(idx, func, file, line) +@device_function function report_exception_frame(idx, func, file, line) if claim_output() != 0 SPIRVIntrinsics.@printf(" [%d] %s at %s:%d\n", idx, func, file, line) end @@ -89,11 +93,13 @@ function report_exception_frame(idx, func, file, line) end # the work-item that threw stops executing after this -function signal_exception() +@device_function function signal_exception() atomic_store_device!(exception_field(Val(:status)), Int32(1)) return end +# these use no device intrinsics, so they can stay host methods. as overlays, inference could +# constant-fold calls to them by running the host method, which throws. malloc(sz) = C_NULL report_oom(sz) = return @@ -107,7 +113,7 @@ struct KernelState exception_info::UInt64 end -@inline @generated kernel_state() = GPUCompiler.kernel_state_value(KernelState) +@device_function @inline @generated kernel_state() = GPUCompiler.kernel_state_value(KernelState) ## intrinsics for adding and accessing additional kernel arguments @@ -132,5 +138,5 @@ end end for name in [:random_keys, :random_counters] - @eval @inline $name() = additional_arg_value(LLVMPtr{UInt32, AS.Workgroup}, Val($(QuoteNode(name)))) + @eval @device_function @inline $name() = additional_arg_value(LLVMPtr{UInt32, AS.Workgroup}, Val($(QuoteNode(name)))) end diff --git a/test/Project.toml b/test/Project.toml index c191d6ced..0136a8024 100644 --- a/test/Project.toml +++ b/test/Project.toml @@ -2,6 +2,7 @@ Adapt = "79e6a3ab-5dfb-504d-930d-738a2a938a0e" Aqua = "4c88cf16-eb10-579e-8560-4a9242c79595" FileCheck = "4e644321-382b-4b05-b0b6-5d23c3d944fb" +GPUToolbox = "096a3bc2-3ced-46d0-87f4-dd12716f4bfc" InteractiveUtils = "b77e0a4c-d291-57a0-90e8-8db25a27a240" KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c" KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf" diff --git a/test/runtests.jl b/test/runtests.jl index 4d09525ca..e1714a1a1 100644 --- a/test/runtests.jl +++ b/test/runtests.jl @@ -1,6 +1,7 @@ using KernelAbstractions using Random using Test +import GPUToolbox include("quality_assurance.jl") include("linenumbers.jl") @@ -559,6 +560,12 @@ end @test occursin("throw_boundserror", out) end +# sysimage builds compile methods for the host, which must not reach device-only code, or +# the sysimage references device intrinsics and fails to load (JuliaGPU/GPUCompiler.jl#611) +@testset "host leaks" begin + @test isempty(GPUToolbox.Overlays.host_leaks(KernelAbstractions)) +end + @testset "POCL compilation cache" begin mod = @eval module $(gensym()) @noinline child() = return From f475f93f8f18badd81bc9001ebaaced98d3ac2d8 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 28 Oct 2026 08:40:57 +0100 Subject: [PATCH 2/2] Precompile a kernel launch on Julia 1.11 on Windows again On Julia 1.11, the launch in the precompile workload left GPUCompiler's runtime in the package image, and on Windows lld links package images eagerly. The device symbols that runtime referenced made the link fail, so the workload skipped the launch there. The POCL runtime no longer reaches device intrinsics on the host, and GPUCompiler 2.19 makes its boxing functions refer to type tags it can link, so drop the gate. Co-Authored-By: Claude Opus 5.5 --- Project.toml | 2 +- src/precompile.jl | 6 +----- test/runtests.jl | 3 +-- 3 files changed, 3 insertions(+), 8 deletions(-) diff --git a/Project.toml b/Project.toml index 5afd14bed..36ba41a4b 100644 --- a/Project.toml +++ b/Project.toml @@ -49,7 +49,7 @@ oneAPIExt = "oneAPI" AMDGPU = "2" Adapt = "0.4, 1.0, 2.0, 3.0, 4" Atomix = "1.5" -GPUCompiler = "2.18" +GPUCompiler = "2.19" GPUToolbox = "3.5" KernelInterface = "0.4" LLVM = "10" diff --git a/src/precompile.jl b/src/precompile.jl index ef7bce19c..cf54ef37c 100644 --- a/src/precompile.jl +++ b/src/precompile.jl @@ -8,15 +8,11 @@ @inbounds A[i] = fma(2, B[i], 1) end -# on Julia 1.11, the launch leaks GPUCompiler's runtime into the package image, which then -# fails to link on Windows -const launch_in_workload = !(Sys.iswindows() && v"1.11-" <= VERSION < v"1.12-") - # whether the workload below launched its kernel, checked by the tests const precompiled_launch = Ref(false) PrecompileTools.@setup_workload begin - if launch_in_workload && POCL.nanoOpenCL.pocl_standalone_jll.is_available() && + if POCL.nanoOpenCL.pocl_standalone_jll.is_available() && POCL.SPIRV_LLVM_Backend_jll.is_available() && POCL.SPIRV_Tools_jll.is_available() try # keep PoCL's kernel cache out of the user's cache directory, and don't have it diff --git a/test/runtests.jl b/test/runtests.jl index e1714a1a1..a6615bd51 100644 --- a/test/runtests.jl +++ b/test/runtests.jl @@ -24,8 +24,7 @@ end @testset "Precompilation" begin # without it, the first kernel launch on the CPU back-end takes many seconds workloads = Base.JLOptions().use_compiled_modules == 1 && - KernelAbstractions.PrecompileTools.workload_enabled(KernelAbstractions) && - KernelAbstractions.launch_in_workload + KernelAbstractions.PrecompileTools.workload_enabled(KernelAbstractions) @test KernelAbstractions.precompiled_launch[] skip = !workloads end