From fa9a42e3df4f64cd8bdb3979b1a21aa819d35eae Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Fri, 21 Aug 2026 21:14:06 +0000 Subject: [PATCH 1/2] Override the InexactError/DomainError constructors on device; add a compile-smoke grid MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit GPUCompiler 2.2.2 (JuliaGPU/GPUCompiler.jl#899, `compilesig_invokes=false`) stopped eliding the `DomainError` allocation inside `Base.Math.exponent`'s local `throw1` closure. `sqrt(::Complex)` reaches it through `ssqs`, and oneAPI.jl provides no device `malloc`, so the gpuarrays/statistics test died with InvalidIRError: unsupported call to an unknown function (call to gpu_malloc) The allocation is the exception object. Base constructs `InexactError` and `DomainError` directly in several places — `exponent`'s closures, `ssqs`'s `Int(::Float)` conversion, `_cpow`, `Int32(::Float32)`, `round(Int, ::Float64)` — none of which goes through a throw helper that `src/device/quirks.jl` could override individually. An exception object only exists to be thrown, so overlay the constructors themselves with `@print_and_throw`; that covers every site at once. The overlays are `@inline` and specialized on purpose: a `@noinline`/`@nospecialize` version has to box its `Float32` argument, which is the same allocation by another route (and the reason `sqrt(::ComplexF32)` never compiled). `test/device/codegen.jl` compiles ~200 (function, eltype) cells through the validating pipeline (`compile_to_obj`; the `code_*` reflection entry points skip validation) without launching anything, so the next "f(::T) stopped compiling" names its cell instead of failing deep inside a GPUArrays test. The cells it still marks broken are SPIR-V code-generator limitations, not allocations: the Khronos translator rejects `_cpow`'s `i63` and `llvm.smul.with.overflow`, the LLVM back-end cannot select `G_SADDO` and mis-legalizes `_cpow`. Verified on an Aurora debug node with GPUCompiler 2.2.2 (LTS stack): device/codegen, execution and gpuarrays/statistics pass (322/6 broken/0). --- src/device/quirks.jl | 17 ++++ test/device/codegen.jl | 204 +++++++++++++++++++++++++++++++++++++++++ 2 files changed, 221 insertions(+) create mode 100644 test/device/codegen.jl diff --git a/src/device/quirks.jl b/src/device/quirks.jl index 987922a5..610b01f8 100644 --- a/src/device/quirks.jl +++ b/src/device/quirks.jl @@ -29,6 +29,23 @@ end @device_override @noinline Core.throw_inexacterror(f::Symbol, ::Type{T}, val) where {T} = @print_and_throw "Inexact conversion" +# Base also constructs these exceptions directly, without a throw helper that could be +# overridden: `Int32(::Float32)` and `round(Int, ::Float64)` (float.jl), `x^y` for `Complex` +# (`_cpow`), and the local `throw1`/`throw2` closures of `exponent` (math.jl) that +# `sqrt(::Complex)` reaches through `ssqs`. Without a device `malloc` the allocation of the +# exception object fails compilation whenever the optimizer does not elide it (GPUCompiler +# 2.2.2 stopped doing so for `exponent`). An exception object only exists to be thrown, so +# replacing the constructors covers every such site at once, at the cost of the specific +# message. Unlike Base's `@nospecialize`d inner constructors these are specialized and +# inlined: a `@noinline` callee taking `Any` would have to box its (e.g. `Float32`) argument, +# which is itself an allocation. +@device_override @inline Core.InexactError(f::Symbol, args...) = + @print_and_throw "Inexact conversion" +@device_override @inline Core.DomainError(val) = + @print_and_throw "Argument outside the domain of the function" +@device_override @inline Core.DomainError(val, msg::AbstractString) = + @print_and_throw "Argument outside the domain of the function" + # abstractarray.jl @device_override @noinline Base.throw_boundserror(A, I) = @print_and_throw "Out-of-bounds array access" diff --git a/test/device/codegen.jl b/test/device/codegen.jl new file mode 100644 index 00000000..c1547a30 --- /dev/null +++ b/test/device/codegen.jl @@ -0,0 +1,204 @@ +# Compile-only coverage of Base/Math functions on the oneAPI target. +# +# oneAPI.jl has no device `malloc`, so any heap allocation that survives optimization — in +# practice an exception object on a throw path that `src/device/quirks.jl` does not cover — +# fails compilation with +# +# InvalidIRError: unsupported call to an unknown function (call to gpu_malloc) +# +# Which throw paths survive depends on what the optimizer happens to delete, so a routine +# Julia/GPUCompiler/LLVM bump can silently change the answer (GPUCompiler 2.2.2 did, for +# `sqrt(::Complex)` via `exponent`). This grid compiles each (function, eltype) cell through +# the real pipeline — the same target, method table and validation a launch uses — without +# launching anything, so a regression names its cell instead of failing three layers deep in +# a GPUArrays test. +# +# Cells known to fail are listed in `BROKEN` with the reason; a cell that starts passing is +# reported as an unexpected pass, which is the signal to remove it from the list. + +import LinearAlgebra + +# `oneAPI.code_llvm` and friends disable IR validation, so go through the validating path +# `zefunction` takes: the job is built for the current device (LTS or rolling back-end, fp +# capabilities, quirks method table) and compiled to SPIR-V, which also exercises the +# translator. Bypasses the kernel cache so cells neither pollute nor hit it. +function compiles(f, tt) + config = oneAPI.compiler_config(device(); kernel = true) + job = oneAPI.CompilerJob(oneAPI.methodinstance(typeof(f), tt), config) + oneAPI.compile_to_obj(job) + return true +end + +# The value flows through memory so the call cannot be folded away. +function smoke_kernel(f, out, x) + @inbounds out[1] = f(x[1]) + return +end + +const DevVec{T} = oneDeviceVector{T, oneAPI.AS.CrossWorkgroup} + +# Result type from a host evaluation where possible; the GPU interpreter otherwise (it +# respects the device overrides, unlike the host). +function result_type(f, T) + R = try + typeof(f(one(T))) + catch + oneAPI.return_type(f, Tuple{T}) + end + # a throwing host evaluation on a `Union{}`-returning cell is still worth compiling + return R === Union{} ? Nothing : R +end + +# Cells that do not compile today, with the reason. None of these is the `gpu_malloc` +# failure; they are limitations of the SPIR-V code generator the active stack uses (the +# Khronos translator on the LTS stack, the LLVM SPIR-V back-end otherwise, see +# `_compiler_config`). Remove an entry once the cell compiles again. +const BROKEN = Set{Tuple{String, DataType}}( + if oneL0.LTS[] + [ + # llvm-spirv: `InvalidBitWidth: 63` — `_cpow` narrows an `Int` to `i63` + ("x^x", ComplexF16), ("x^x", ComplexF32), ("x^x", ComplexF64), + # llvm-spirv: `Unexpected llvm intrinsic: llvm.smul.with.overflow` + ("checked_mul", Int32), ("checked_mul", Int64), + ] + else + [ + # SPIR-V back-end: malformed `select`/`phi` after its own legalization of `_cpow` + ("x^x", ComplexF16), ("x^x", ComplexF32), + # SPIR-V back-end: `cannot select: G_SADDO` + ("checked_add", Int32), ("checked_add", Int64), + ] + end +) + +function smoke(label, f, T) + R = result_type(f, T) + tt = Tuple{typeof(f), DevVec{R}, DevVec{T}} + broken = (label, T) in BROKEN + @testset "$label($T)" begin + if broken + # the SPIR-V tools print their diagnostics to stderr; known failures stay quiet + redirect_stderr(devnull) do + @test compiles(smoke_kernel, tt) broken = true + end + else + @test compiles(smoke_kernel, tt) + end + end + return +end + +# Named wrappers: the obvious spelling would take a different Base path (`x^-1` is +# `literal_pow`/`inv`, never `throw_domerr_powbysq`) or needs a second operand. +negpow(x) = x^(-one(x)) +powsame(x) = x^x +divsame(x) = x / x +intdiv(x) = div(x, x) +intrem(x) = rem(x, x) +intfld(x) = fld(x, x) +intcld(x) = cld(x, x) +intmod(x) = mod(x, x) +hypotsame(x) = hypot(x, x) +atan2same(x) = atan(x, x) +copysignsame(x) = copysign(x, -x) +checked_add_same(x) = Base.checked_add(x, x) +checked_sub_same(x) = Base.checked_sub(x, x) +checked_mul_same(x) = Base.checked_mul(x, x) +to_int32(x) = Int32(x) +to_int64(x) = Int64(x) +trunc_int32(x) = trunc(Int32, x) +round_int64(x) = round(Int64, x) +to_complexf32(x) = ComplexF32(x) + +float_types = DataType[Float32] +float64_supported && push!(float_types, Float64) +float16_supported && push!(float_types, Float16) +complex_types = DataType[Complex{T} for T in float_types] +int_types = DataType[Int32, Int64, UInt32, UInt64] + +@testset "unary real" begin + for (label, f) in ( + ("sqrt", sqrt), ("cbrt", cbrt), + ("exp", exp), ("exp2", exp2), ("exp10", exp10), ("expm1", expm1), + ("log", log), ("log2", log2), ("log10", log10), ("log1p", log1p), + ("sin", sin), ("cos", cos), ("tan", tan), + ("asin", asin), ("acos", acos), ("atan", atan), + ("sinh", sinh), ("cosh", cosh), ("tanh", tanh), + ("abs", abs), ("abs2", abs2), ("sign", sign), ("inv", inv), + ("exponent", exponent), ("significand", significand), ("frexp", frexp), + ("trunc", trunc), ("round", round), ("floor", floor), ("ceil", ceil), + ), T in float_types + smoke(label, f, T) + end +end + +@testset "unary complex" begin + for (label, f) in ( + ("sqrt", sqrt), ("exp", exp), ("log", log), ("sin", sin), ("cos", cos), + ("abs", abs), ("abs2", abs2), ("sign", sign), ("inv", inv), ("angle", angle), + ), T in complex_types + smoke(label, f, T) + end +end + +@testset "binary real" begin + for (label, f) in ( + ("x^x", powsame), ("x/x", divsame), ("hypot", hypotsame), ("atan2", atan2same), + ("mod", intmod), ("rem", intrem), ("div", intdiv), ("fld", intfld), ("cld", intcld), + ("copysign", copysignsame), + ), T in float_types + smoke(label, f, T) + end +end + +@testset "binary complex" begin + for (label, f) in (("x^x", powsame), ("x/x", divsame)), T in complex_types + smoke(label, f, T) + end +end + +@testset "integer" begin + for (label, f) in ( + ("x^-1", negpow), + ("div", intdiv), ("rem", intrem), ("fld", intfld), ("mod", intmod), + ("checked_add", checked_add_same), ("checked_sub", checked_sub_same), + ("checked_mul", checked_mul_same), + ), T in int_types + smoke(label, f, T) + end +end + +@testset "conversions" begin + for (label, f, T) in ( + ("Int32", to_int32, Float32), ("Int64", to_int64, Float64), + ("trunc(Int32)", trunc_int32, Float32), ("round(Int64)", round_int64, Float64), + ("ComplexF32", to_complexf32, ComplexF64), + ) + (T == Float64 || T == ComplexF64) && !float64_supported && continue + smoke(label, f, T) + end +end + +# Bounds-checked array access (the `throw_boundserror` quirk) and the `Diagonal` +# `setindex!` quirk; compile only, never launched. +function checked_index_kernel(out, x) + out[2] = x[2] + return +end +function diagonal_setindex_kernel(out, x) + D = LinearAlgebra.Diagonal(x) + D[1, 2] = x[1] + out[1] = D[1, 1] + return +end + +@testset "array" begin + for T in (Int32, Float32) + @testset "checked indexing($T)" begin + @test compiles(checked_index_kernel, Tuple{DevVec{T}, DevVec{T}}) + end + @testset "Diagonal setindex!($T)" begin + @test compiles(diagonal_setindex_kernel, Tuple{DevVec{T}, DevVec{T}}) + end + end +end From a9cd2880426ca00deafa6b5026582e36aadb5c9a Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Fri, 21 Aug 2026 22:15:25 +0000 Subject: [PATCH 2/2] Add a device runtime: host-visible exceptions and a per-work-item heap for malloc MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Device exceptions used to be silent. `signal_exception` was a no-op and the SPIR-V target lowers the trap behind it to a plain return, so a kernel that threw — a bounds error, a domain error, anything — just stopped that work-item and the host never heard about it. And anything that allocated a Julia object on the device did not compile at all: GPUCompiler lowers heap allocations to a back-end `malloc` that oneAPI.jl never provided, so every throw path whose exception object the optimizer failed to delete ended in InvalidIRError: unsupported call to an unknown function (call to gpu_malloc) which is what GPUCompiler 2.2.2 exposed for `sqrt(::Complex)`. Both are now served through a kernel state, the hidden first argument GPUCompiler threads to every device function: - `exception_flag` points at a 16-byte host USM buffer per (context, device). `signal_exception` and `report_oom` set its words; the host reads and atomically clears them whenever it synchronizes a stream — `synchronize()`, `@sync`, copying back to the host — and raises a `KernelException` (with a note when the heap ran out). The check is deliberately absent from `synchronize_all_streams`, which runs from finalizers. The state is prepended in the generated `call`, not in `onecall`, which stays a raw launch primitive for foreign SPIR-V. - `heap` points at a bump arena in the work-item's private memory, and `malloc` bump-allocates from it without atomics. A global-memory heap was tried first and does not work: Julia's boxed objects live in address space 0 once GPUCompiler strips its address spaces, which is private memory to SPIR-V and to Intel's compiler, and stores through such a pointer into a global buffer are silently lost (verified on a Max 1550: allocations happened, the values never arrived). SPIR-V has no per-invocation module-scope storage either (the translator rejects address-space-0 globals), so `add_private_heap!` allocas the arena in the kernel entry after the kernel-state passes and patches its pointer into the state the entry hands on. Only kernels whose code reaches `gpu_malloc` carry the arena; the heap is fresh on every launch and nothing is freed. Exhausting the `PRIVATE_HEAP_SIZE` (1 KiB) is reported as an out-of-memory `KernelException`, never a silent failure. The constructor overlays in quirks.jl stay: they print a reason, which the heap path does not, and keep allocations off the hot error paths. Verified on an Aurora debug node (LTS stack, GPUCompiler 2.2.2): boxed values round-trip through the heap, 4096 concurrent work-items allocate correctly, exhaustion is reported, and the full test suite passes (12830 pass, 61 broken; the one failure, broadcast Float16, is the known host AVX512-FP16 issue and passes with the CI's JIT target). --- docs/src/kernels.md | 28 +++++++++++ src/compiler/compilation.jl | 40 ++++++++++++++++ src/compiler/execution.jl | 17 ++++++- src/context.jl | 4 ++ src/device/quirks.jl | 18 +++---- src/device/runtime.jl | 64 +++++++++++++++++++++++-- src/exceptions.jl | 83 ++++++++++++++++++++++++++++++++ src/oneAPI.jl | 1 + test/device/codegen.jl | 19 +++----- test/execution.jl | 96 +++++++++++++++++++++++++++++++++++++ 10 files changed, 342 insertions(+), 28 deletions(-) create mode 100644 src/exceptions.jl diff --git a/docs/src/kernels.md b/docs/src/kernels.md index 07e50c7a..41437303 100644 --- a/docs/src/kernels.md +++ b/docs/src/kernels.md @@ -62,3 +62,31 @@ and correspond to the standard OpenCL built-in functions. Note that the indices are 1-based, so they can be used to index Julia arrays directly. See [Device Intrinsics](device.md) for the full list. + +## Exceptions and Dynamic Allocation + +Kernels can throw. An exception on the device aborts the work-item that threw it and is +reported on the host as a `KernelException` at the next synchronization — `synchronize()`, +`oneAPI.@sync`, or copying data back with `Array` — after the reason was printed by the +device: + +```julia +julia> function kernel(a) + a[2] = 1f0 # bounds-checked + return + end; + +julia> @oneapi kernel(oneArray(Float32[0])); + +julia> synchronize() +ERROR: Out-of-bounds array access. +ERROR: KernelException: exception thrown during kernel execution on device Intel(R) Data Center GPU Max 1550 +``` + +Exception objects that survive optimization, and any other Julia object that has to be +heap-allocated inside a kernel (for example a `Ref` passed to a `@noinline` function), are +allocated from a small per-work-item heap in private memory; objects never outlive the +work-item that created them and are never freed. The heap is limited to +`oneAPI.PRIVATE_HEAP_SIZE` bytes per work-item, and exhausting it is reported as a +`KernelException` as well, rather than failing silently. Code on a hot path should not +allocate. diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl index a8c117a8..9554b765 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -30,6 +30,10 @@ end GPUCompiler.runtime_module(::oneAPICompilerJob) = oneAPI +# hidden first kernel argument carrying the exception flag and the private heap; see +# src/device/runtime.jl, src/exceptions.jl and `add_private_heap!` below +GPUCompiler.kernel_state_type(::oneAPICompilerJob) = KernelState + GPUCompiler.method_table_view(job::oneAPICompilerJob) = GPUCompiler.StackedMethodTable(job.world, method_table, SPIRVIntrinsics.method_table) @@ -64,6 +68,9 @@ end # finish_ir! runs later in the pipeline, after optimizations that create nested insertvalue function GPUCompiler.finish_ir!(job::oneAPICompilerJob, mod::LLVM.Module, entry::LLVM.Function) + # before the kernel state is turned into a by-reference argument below + job.config.kernel && add_private_heap!(mod, entry) + entry = invoke(GPUCompiler.finish_ir!, Tuple{CompilerJob{SPIRVCompilerTarget}, typeof(mod), typeof(entry)}, job, mod, entry) @@ -92,6 +99,39 @@ function GPUCompiler.finish_ir!(job::oneAPICompilerJob, mod::LLVM.Module, return entry end +# Give the kernel a private heap for `malloc` (src/device/runtime.jl): allocate the arena in +# the entry block of the kernel, initialize its cursor, and thread the pointer into the kernel +# state that the entry passes on to every device function. Runs after the kernel-state passes, +# so the state is the entry's first (by-value) parameter; only kernels whose code reaches +# `gpu_malloc` pay for the arena. +function add_private_heap!(mod::LLVM.Module, entry::LLVM.Function) + haskey(functions(mod), "gpu_malloc") || return false + T_state = convert(LLVMType, KernelState) + params = parameters(entry) + (isempty(params) || value_type(params[1]) != T_state) && return false + state = params[1] + isempty(uses(state)) && return false + + users = LLVM.Value[user(use) for use in uses(state)] + @dispose builder = IRBuilder() begin + position!(builder, first(instructions(first(blocks(entry))))) + T_i8 = LLVM.Int8Type() + T_i32 = LLVM.Int32Type() + arena = alloca!(builder, LLVM.ArrayType(T_i8, HEAP_HEADER + PRIVATE_HEAP_SIZE), "private_heap") + alignment!(arena, 16) + store!(builder, ConstantInt(T_i32, 0), arena) + heap_field = findfirst(==(:heap), fieldnames(KernelState)) - 1 + new_state = insert_value!(builder, state, arena, heap_field, "state_with_heap") + for u in users + ops = operands(u) + for i in 1:length(ops) + ops[i] == state && (ops[i] = new_state) + end + end + end + return true +end + # Flatten nested insertvalue instructions # This works around a bug in Intel's SPIR-V runtime where OpCompositeInsert # with nested array indices corrupts adjacent struct fields. diff --git a/src/compiler/execution.jl b/src/compiler/execution.jl index bd874d67..f1f387e1 100644 --- a/src/compiler/execution.jl +++ b/src/compiler/execution.jl @@ -195,6 +195,12 @@ abstract type AbstractKernel{F,TT} end end end + # the kernel state is the hidden first argument of every compiled kernel (see + # `GPUCompiler.kernel_state_type`); `onecall` itself stays agnostic so that it can + # also launch kernels that were not compiled by us + pushfirst!(call_t, KernelState) + pushfirst!(call_args, :(kernel.state)) + # finalize types call_tt = Base.to_tuple_type(call_t) @@ -209,6 +215,8 @@ end struct HostKernel{F,TT} <: AbstractKernel{F,TT} f::F fun::ZeKernel + # for the context and device `fun` was linked against + state::KernelState end # Upper bound on the spill (scratch) memory a single work-group may require, in bytes. @@ -297,7 +305,7 @@ function zefunction(f::F, tt::TT=Tuple{}; kwargs...) where {F,TT} # about world age here, as GPUCompiler already does and will return a different object h = hash(fun, hash(f, hash(tt))) get!(_kernel_instances, h) do - HostKernel{F,tt}(f, fun) + HostKernel{F, tt}(f, fun, kernel_state(ctx, dev)) end::HostKernel{F,tt} end end @@ -346,7 +354,10 @@ end spill > s.scratch_hwm && scratch_hedge!(s, spill) append_launch!(s.list, kernel, groups) - oneL0.sync_each_submission() && oneL0.synchronize(s.list) + if oneL0.sync_each_submission() + oneL0.synchronize(s.list) + check_exceptions(s.ctx, s.dev) + end return end @@ -358,6 +369,8 @@ end execute!(queue) do list append_launch!(list, kernel, groups) end + oneL0.sync_each_submission() && check_exceptions(queue.context, queue.device) + return end # Slow path of the scratch hedge, firing once per (stream, spill tier): retire in-flight diff --git a/src/context.jl b/src/context.jl index 409a9d7a..bf0d9fb4 100644 --- a/src/context.jl +++ b/src/context.jl @@ -405,6 +405,10 @@ function oneL0.synchronize(s::oneStream) oneL0.synchronize(q) s.mkl_dirty = false end + # every user-facing synchronization funnels through here, so this is where a device + # exception surfaces (src/exceptions.jl); `synchronize_all_streams` deliberately does + # not check, it runs from finalizers + check_exceptions(s.ctx, s.dev) return end diff --git a/src/device/quirks.jl b/src/device/quirks.jl index 610b01f8..27e0968a 100644 --- a/src/device/quirks.jl +++ b/src/device/quirks.jl @@ -32,13 +32,12 @@ end # Base also constructs these exceptions directly, without a throw helper that could be # overridden: `Int32(::Float32)` and `round(Int, ::Float64)` (float.jl), `x^y` for `Complex` # (`_cpow`), and the local `throw1`/`throw2` closures of `exponent` (math.jl) that -# `sqrt(::Complex)` reaches through `ssqs`. Without a device `malloc` the allocation of the -# exception object fails compilation whenever the optimizer does not elide it (GPUCompiler -# 2.2.2 stopped doing so for `exponent`). An exception object only exists to be thrown, so -# replacing the constructors covers every such site at once, at the cost of the specific -# message. Unlike Base's `@nospecialize`d inner constructors these are specialized and -# inlined: a `@noinline` callee taking `Any` would have to box its (e.g. `Float32`) argument, -# which is itself an allocation. +# `sqrt(::Complex)` reaches through `ssqs`. Left alone, such a throw allocates the exception +# object on the device heap and signals the host without printing anything; replacing the +# constructors covers every such site at once with a printed reason (at the cost of its +# specificity) and keeps the heap out of it. Unlike Base's `@nospecialize`d inner +# constructors these are specialized and inlined: a `@noinline` callee taking `Any` would +# have to box its (e.g. `Float32`) argument, which is itself an allocation. @device_override @inline Core.InexactError(f::Symbol, args...) = @print_and_throw "Inexact conversion" @device_override @inline Core.DomainError(val) = @@ -55,7 +54,8 @@ end @print_and_throw "sincos(x) is only defined for finite x." # diagonal.jl -# XXX: remove when we have malloc +# Base's version throws an ArgumentError built from a string; this one prints and keeps the +# device heap out of the hot path. import LinearAlgebra @device_override function Base.setindex!(D::LinearAlgebra.Diagonal, v, i::Int, j::Int) @boundscheck checkbounds(D, i, j) @@ -68,7 +68,7 @@ import LinearAlgebra end # number.jl -# XXX: remove when we have malloc +# Base's version throws a BoundsError; same reasoning as above. @device_override @inline function Base.getindex(x::Number, I::Integer...) @boundscheck all(isone, I) || @print_and_throw "Out-of-bounds access of scalar value" diff --git a/src/device/runtime.jl b/src/device/runtime.jl index 2f48b0f1..508ac31c 100644 --- a/src/device/runtime.jl +++ b/src/device/runtime.jl @@ -1,9 +1,50 @@ # device runtime libraries +# +# GPUCompiler resolves the back-end runtime by name in this module (`runtime_module`): +# `signal_exception`, `report_*` and `malloc` below are compiled into the runtime library +# that is linked into every kernel, with `malloc` becoming the `gpu_malloc` symbol that +# `gc_pool_alloc` — and so every heap allocation that survives optimization — calls. -## Julia library +## kernel state + +# Passed by value as the hidden first argument of every kernel and threaded by GPUCompiler +# to every device function that calls `kernel_state()`. The exception flag is filled in on +# the host (`kernel_state` in src/exceptions.jl, once per linked kernel); the heap pointer +# is patched in on the device, in the kernel's entry block (`add_private_heap!` in +# src/compiler/compilation.jl), because it points at private memory. +struct KernelState + # host USM, 16 bytes, read by the host after synchronization and cleared by it: + # [1] nonzero after `signal_exception`, [2] nonzero after `report_oom` + exception_flag::LLVMPtr{Int32, AS.CrossWorkgroup} + # per-work-item private memory: a `HEAP_HEADER`-byte header whose first word is the + # bump cursor, followed by `PRIVATE_HEAP_SIZE` allocatable bytes; null in kernels that + # do not allocate + heap::Ptr{UInt8} +end + +# Bytes of private memory set aside for dynamic allocations, per work-item. Julia objects +# allocated in a kernel never outlive the work-item that created them — exception objects +# on a throw path, boxes handed to a `@noinline` callee — so private memory is the right +# place for them, and the only one: address space 0, which Julia's boxed objects live in +# after GPUCompiler strips its address spaces, is private memory to SPIR-V and Intel's +# compiler, and a store through an address-space-0 pointer into global memory is silently +# lost. The arena is only materialized in kernels whose code calls `malloc`. +const PRIVATE_HEAP_SIZE = 1024 +const HEAP_HEADER = 16 # keeps the first allocation 16-byte aligned + +@inline @generated kernel_state() = GPUCompiler.kernel_state_value(KernelState) + + +## exceptions function signal_exception() + unsafe_store!(kernel_state().exception_flag, Int32(1), 1) + return +end + +function report_oom(sz) + unsafe_store!(kernel_state().exception_flag, Int32(1), 2) return end @@ -15,8 +56,6 @@ function report_exception(ex) return end -report_oom(sz) = return #@cuprintf("ERROR: Out of dynamic GPU memory (trying to allocate %i bytes)\n", sz) - function report_exception_name(ex) # @cuprintf(""" # ERROR: a %s was thrown during kernel execution. @@ -31,6 +70,21 @@ function report_exception_frame(idx, func, file, line) end -## SPIRV libraries +## dynamic memory allocation -# TODO +# Bump allocator over the work-item's private arena: nothing is ever freed (the arena dies +# with the work-item), and exhaustion returns null, which `gc_pool_alloc` turns into +# `report_oom` + `OutOfMemoryError`, i.e. a loud `KernelException` on the host rather than +# a silent failure. The cursor is private to the work-item, so no atomics are needed. +function malloc(sz::Csize_t) + heap = kernel_state().heap + heap == C_NULL && return C_NULL + sz > PRIVATE_HEAP_SIZE && return C_NULL + bytes = (UInt32(sz) + UInt32(15)) & ~UInt32(15) + cursor = convert(Ptr{UInt32}, heap) + old = unsafe_load(cursor) + new = old + bytes + new > PRIVATE_HEAP_SIZE && return C_NULL + unsafe_store!(cursor, new) + return Ptr{Cvoid}(heap + HEAP_HEADER + old) +end diff --git a/src/exceptions.jl b/src/exceptions.jl new file mode 100644 index 00000000..5bf733b5 --- /dev/null +++ b/src/exceptions.jl @@ -0,0 +1,83 @@ +# device exceptions +# +# Every kernel receives a `KernelState` (src/device/runtime.jl) that names a host-visible +# flag owned by the (context, device) it was linked for. The device runtime sets the flag +# when a kernel throws; the host reads it whenever it synchronizes a stream and raises a +# `KernelException`. + +export KernelException + + +## exception type + +""" + KernelException + +Thrown on the host, at the next synchronization, after a kernel threw an exception on the +device. The device runtime prints the reason (e.g. `ERROR: Out-of-bounds array access.`) +to the standard output of the process at the time of the throw; this exception only carries +the device and whether the work-item ran out of private heap memory along the way. +""" +struct KernelException <: Exception + dev::ZeDevice + oom::Bool +end + +function Base.showerror(io::IO, err::KernelException) + name = oneL0.properties(err.dev).name + print(io, "KernelException: exception thrown during kernel execution on device $name") + if err.oom + print( + io, " (a work-item allocated more than the $(PRIVATE_HEAP_SIZE) bytes of ", + "private heap memory available to it)" + ) + end + return +end + + +## exception flags + +# one 16-byte host USM buffer per (context, device): [1] exception, [2] oom (Int32 each) +const exception_flags = Dict{Tuple{ZeContext, ZeDevice}, oneL0.HostBuffer}() +const exception_flags_lock = ReentrantLock() + +function exception_flag(ctx::ZeContext, dev::ZeDevice) + return Base.@lock exception_flags_lock get!(exception_flags, (ctx, dev)) do + flag = oneL0.host_alloc(ctx, 16, 16) + # a pointer embedded in the kernel state is an indirect access as far as Level Zero + # is concerned, so the buffer needs explicit residency + oneL0.make_resident(ctx, dev, flag) + p = convert(Ptr{Int32}, flag) + unsafe_store!(p, Int32(0), 1) + unsafe_store!(p, Int32(0), 2) + flag + end +end + +# the kernel state for kernels linked against `ctx`/`dev`; built once per `HostKernel`. +# The private heap pointer is filled in on the device (`add_private_heap!`). +function kernel_state(ctx::ZeContext, dev::ZeDevice) + flag = exception_flag(ctx, dev) + return KernelState( + reinterpret(LLVMPtr{Int32, AS.CrossWorkgroup}, pointer(flag)), + C_NULL + ) +end + + +## host-side check + +# Called after a stream has been synchronized. Clears the flag words so the exception is +# reported once; the clear is an atomic swap so two tasks synchronizing the same device +# cannot both report one throw. +function check_exceptions(ctx::ZeContext, dev::ZeDevice) + flag = Base.@lock exception_flags_lock get(exception_flags, (ctx, dev), nothing) + flag === nothing && return + p = convert(Ptr{Int32}, flag) + unsafe_load(p, 1) == 0 && return + thrown = Core.Intrinsics.atomic_pointerswap(p, Int32(0), :sequentially_consistent) + oom = Core.Intrinsics.atomic_pointerswap(p + sizeof(Int32), Int32(0), :sequentially_consistent) + thrown == 0 && return + throw(KernelException(dev, oom != 0)) +end diff --git a/src/oneAPI.jl b/src/oneAPI.jl index 7baebfa4..53dc5212 100644 --- a/src/oneAPI.jl +++ b/src/oneAPI.jl @@ -51,6 +51,7 @@ include("context.jl") include("memory.jl") include("pool.jl") include("array.jl") +include("exceptions.jl") # compiler implementation include("compiler/compilation.jl") diff --git a/test/device/codegen.jl b/test/device/codegen.jl index c1547a30..2a9c2a5e 100644 --- a/test/device/codegen.jl +++ b/test/device/codegen.jl @@ -1,17 +1,12 @@ # Compile-only coverage of Base/Math functions on the oneAPI target. # -# oneAPI.jl has no device `malloc`, so any heap allocation that survives optimization — in -# practice an exception object on a throw path that `src/device/quirks.jl` does not cover — -# fails compilation with -# -# InvalidIRError: unsupported call to an unknown function (call to gpu_malloc) -# -# Which throw paths survive depends on what the optimizer happens to delete, so a routine -# Julia/GPUCompiler/LLVM bump can silently change the answer (GPUCompiler 2.2.2 did, for -# `sqrt(::Complex)` via `exponent`). This grid compiles each (function, eltype) cell through -# the real pipeline — the same target, method table and validation a launch uses — without -# launching anything, so a regression names its cell instead of failing three layers deep in -# a GPUArrays test. +# Whether a Base function compiles for the device depends on which of its throw paths survive +# optimization and whether the device runtime can serve them, so a routine Julia/GPUCompiler/ +# LLVM bump can silently change the answer (GPUCompiler 2.2.2 did, for `sqrt(::Complex)` via +# `exponent`, back when oneAPI.jl had no device `malloc`). This grid compiles each (function, +# eltype) cell through the real pipeline — the same target, method table and validation a +# launch uses — without launching anything, so a regression names its cell instead of failing +# three layers deep in a GPUArrays test. # # Cells known to fail are listed in `BROKEN` with the reason; a cell that starts passing is # reported as an unexpected pass, which is the signal to remove it from the list. diff --git a/test/execution.jl b/test/execution.jl index 1f44c128..a3aa6af3 100644 --- a/test/execution.jl +++ b/test/execution.jl @@ -742,3 +742,99 @@ end end for _ in 1:2]) @test all(results) end + +############################################################################################ + +# Device exceptions reach the host through the kernel state (src/exceptions.jl), and the +# per-work-item private heap behind `malloc` (src/device/runtime.jl) serves the allocations +# that survive optimization. Top-level definitions: a closure capture would not be a bitstype. + +struct ExceptionTestError <: Exception + val::Any +end +# the exception object survives lowering: its `Any` field boxes the Float32 (the +# `DomainError(x, msg)` shape), so it is allocated on the device heap before the throw +@noinline exception_test_thrower(x::Float32) = throw(ExceptionTestError(x)) +@noinline exception_test_consume(r::Base.RefValue{Float32}) = r[] + 1.0f0 + +@testset "device exceptions" begin + # a quirked throw prints the reason and signals; no allocation involved + function boundserror_kernel(a) + a[2] = 1.0f0 + return + end + a = oneArray(Float32[0]) + _, out = @grab_output begin + @oneapi boundserror_kernel(a) + @test_throws KernelException synchronize() + end + @test occursin("Out-of-bounds array access", out) + # reported once: the flag is cleared + synchronize() + + # an un-quirked throw whose exception object lives on the device heap + function alloc_throw_kernel(a) + x = a[1] + x == 0 && exception_test_thrower(x) + a[1] = 2 + return + end + a = oneArray(Float32[0]) + @oneapi alloc_throw_kernel(a) + @test_throws KernelException synchronize() + a = oneArray(Float32[1]) + @oneapi alloc_throw_kernel(a) + @test Array(a) == [2] + + # every user-facing synchronization surfaces it + a = oneArray(Float32[0]) + @test_throws KernelException oneAPI.@sync @oneapi alloc_throw_kernel(a) + a = oneArray(Float32[0]) + @oneapi alloc_throw_kernel(a) + @test_throws KernelException Array(a) +end + +@testset "device heap" begin + # a box handed to a @noinline callee round-trips through the work-item's private heap + function boxing_kernel(a) + i = get_global_id() + a[i] = exception_test_consume(Ref(a[i])) + return + end + a = oneArray(Float32[41]) + @oneapi boxing_kernel(a) + @test Array(a) == [42] + n = 4096 + a = oneArray(Float32.(1:n)) + @oneapi items = 256 groups = n ÷ 256 boxing_kernel(a) + @test Array(a) == Float32.(2:(n + 1)) + + # nothing is freed: allocating more than PRIVATE_HEAP_SIZE bytes in one work-item is a + # KernelException that reports the exhaustion, not a silent failure + function boxing_loop_kernel(a, n) + i = get_global_id() + x = a[i] + for _ in 1:n + x = exception_test_consume(Ref(x)) + end + a[i] = x + return + end + fits = oneAPI.PRIVATE_HEAP_SIZE ÷ 16 - 1 + a = oneArray(Float32.(1:256)) + @oneapi items = 256 boxing_loop_kernel(a, fits) + @test Array(a) == Float32.(1:256) .+ fits + a = oneArray(Float32.(1:256)) + @oneapi items = 256 boxing_loop_kernel(a, 2 * fits) + err = try + synchronize() + nothing + catch e + e + end + @test err isa KernelException && err.oom + # the next launch starts from a fresh heap + a = oneArray(Float32[41]) + @oneapi boxing_kernel(a) + @test Array(a) == [42] +end