From 1be6459e4384ffe5a3d022b181f488cfc1aa455a Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 12 Aug 2026 10:15:35 +0200 Subject: [PATCH 1/3] Serialize mutable Level Zero kernel launches MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ZeKernel argument slots and group size are mutable state shared by cached HostKernel instances. Concurrent launches could interleave their zeKernelSetArgumentValue and zeKernelSetGroupSize calls, causing either command list to capture mixed arguments or the other launch’s group size. Give each kernel its own lock and hold it continuously from the first mutation through append_launch!, while leaving independent kernel handles concurrent. --- lib/level-zero/module.jl | 7 ++++++- src/compiler/execution.jl | 14 ++++++++------ 2 files changed, 14 insertions(+), 7 deletions(-) diff --git a/lib/level-zero/module.jl b/lib/level-zero/module.jl index 0c9ddbb1..68a5af61 100644 --- a/lib/level-zero/module.jl +++ b/lib/level-zero/module.jl @@ -79,6 +79,7 @@ export ZeKernel mutable struct ZeKernel mod::ZeModule handle::ze_kernel_handle_t + lock::ReentrantLock function ZeKernel(mod, name) GC.@preserve name begin @@ -86,7 +87,7 @@ mutable struct ZeKernel handle_ref = Ref{ze_kernel_handle_t}() zeKernelCreate(mod, desc_ref, handle_ref) end - obj = new(mod, handle_ref[]) + obj = new(mod, handle_ref[], ReentrantLock()) finalizer(obj) do obj zeKernelDestroy(obj) @@ -95,6 +96,10 @@ mutable struct ZeKernel end end +Base.lock(kernel::ZeKernel) = lock(getfield(kernel, :lock)) +Base.lock(f::Function, kernel::ZeKernel) = lock(f, getfield(kernel, :lock)) +Base.unlock(kernel::ZeKernel) = unlock(getfield(kernel, :lock)) + Base.unsafe_convert(::Type{ze_kernel_handle_t}, kernel::ZeKernel) = kernel.handle Base.:(==)(a::ZeKernel, b::ZeKernel) = a.handle == b.handle diff --git a/src/compiler/execution.jl b/src/compiler/execution.jl index 327cf8c7..7cb37cff 100644 --- a/src/compiler/execution.jl +++ b/src/compiler/execution.jl @@ -323,13 +323,15 @@ const _kernel_instances = Dict{UInt, Any}() @inline function onecall(kernel::ZeKernel, tt, args...; groups::ZeDim=1, items::ZeDim=1, queue::ZeCommandQueue=global_queue(context(), device())) - for (i, arg) in enumerate(args) - oneL0.arguments(kernel)[i] = arg - end + Base.@lock kernel begin + for (i, arg) in enumerate(args) + oneL0.arguments(kernel)[i] = arg + end - groupsize!(kernel, items) - execute!(queue) do list - append_launch!(list, kernel, groups) + groupsize!(kernel, items) + execute!(queue) do list + append_launch!(list, kernel, groups) + end end end From 5066a72deccdfcf0f0c45d2b7aa9ea6820d75055 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 12 Aug 2026 10:16:14 +0200 Subject: [PATCH 2/3] Protect process-wide context and compiler caches The global context and compiler-configuration dictionaries used unsynchronized get! and check-then-insert paths. Concurrent cold calls could mutate a Dict simultaneously or construct duplicate handles and configurations. Protect each cache with a dedicated lock and keep lookup, construction, and insertion in one atomic get! transaction. --- src/compiler/compilation.jl | 10 +++++----- src/context.jl | 8 ++++++-- 2 files changed, 11 insertions(+), 7 deletions(-) diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl index b49d17b2..a8c117a8 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -302,14 +302,14 @@ end # cache of compiler configurations, per device (but additionally configurable via kwargs) const _toolchain = Ref{Any}() const _compiler_configs = Dict{UInt, oneAPICompilerConfig}() +const compiler_config_lock = ReentrantLock() function compiler_config(dev; kwargs...) h = hash(dev.driver, hash(dev, hash(kwargs))) - config = get(_compiler_configs, h, nothing) - if config === nothing - config = _compiler_config(dev; kwargs...) - _compiler_configs[h] = config + Base.@lock compiler_config_lock begin + get!(_compiler_configs, h) do + _compiler_config(dev; kwargs...) + end end - return config end # Whether the driver's SPIR-V runtime accepts the SPV_KHR_bfloat16 extension. function _driver_supports_bfloat16_spirv(dev=device()) diff --git a/src/context.jl b/src/context.jl index f625a40d..15a1037a 100644 --- a/src/context.jl +++ b/src/context.jl @@ -146,6 +146,7 @@ function is_integrated(dev::ZeDevice=device()) end const global_contexts = Dict{ZeDriver,ZeContext}() +const global_contexts_lock = ReentrantLock() """ context() -> ZeContext @@ -166,8 +167,11 @@ See also: [`context!`](@ref), [`driver`](@ref) """ function context() get!(task_local_storage(), :ZeContext) do - get!(global_contexts, driver()) do - ZeContext(driver()) + drv = driver() + Base.@lock global_contexts_lock begin + get!(global_contexts, drv) do + ZeContext(drv) + end end end end From 1cab516ad25c3ccd9992c5caf916302873c1496d Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 12 Aug 2026 10:17:06 +0200 Subject: [PATCH 3/3] Make the default GPUArrays RNG task-local The process-wide device cache returned the same mutable RNG and state array to tasks that use independent queues. Concurrent seed! and generation calls could therefore update shared device state without queue ordering. Move the per-device RNG dictionary into task-local storage so repeated calls within one task reuse state while separate tasks receive independent RNGs. --- src/gpuarrays.jl | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/src/gpuarrays.jl b/src/gpuarrays.jl index e18b3c72..a0a38167 100644 --- a/src/gpuarrays.jl +++ b/src/gpuarrays.jl @@ -1,9 +1,11 @@ # GPUArrays.jl interface -const GLOBAL_RNGs = Dict{ZeDevice,GPUArrays.RNG}() function GPUArrays.default_rng(::Type{<:oneArray}) dev = device() - get!(GLOBAL_RNGs, dev) do + rngs = get!(task_local_storage(), :oneAPI_GLOBAL_RNGs) do + Dict{ZeDevice,GPUArrays.RNG}() + end + get!(rngs, dev) do N = oneL0.compute_properties(dev).maxTotalGroupSize state = oneArray{NTuple{4, UInt32}}(undef, N) rng = GPUArrays.RNG(state)