From dd2e41fe529cae3d55a3b3f3d8f596df35cd949d Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Mon, 19 Oct 2026 19:02:42 +0200 Subject: [PATCH] Raise a clear error when passing CPU arrays to kernels Mixing a oneArray with an Array in a broadcast, or passing an Array to `@oneapi`, only failed in the compiler's argument validation, which dumps the full argument type and its non-isbits fields. Reject Arrays during argument conversion instead, with a message pointing at `oneArray`. --- src/compiler/execution.jl | 6 ++++++ test/execution.jl | 11 +++++++++++ 2 files changed, 17 insertions(+) diff --git a/src/compiler/execution.jl b/src/compiler/execution.jl index 1ff3523f..92254e7f 100644 --- a/src/compiler/execution.jl +++ b/src/compiler/execution.jl @@ -169,6 +169,12 @@ function Adapt.adapt_storage(to::KernelAdaptor, xs::oneArray{T,N}) where {T,N} oneDeviceArray{T,N,AS.CrossWorkgroup}(size(xs), ptr, xs.maxsize - xs.offset) end +# CPU arrays would otherwise only be rejected by the compiler, with an error that dumps the +# full argument type +Adapt.adapt_storage(::KernelAdaptor, xs::Array) = + throw(ArgumentError("""cannot use a CPU array ($(typeof(xs))) in a oneAPI kernel. + Convert it to a GPU array first, e.g. using `oneArray(x)`.""")) + # Base.RefValue isn't GPU compatible, so provide a compatible alternative. # TODO: port improvements from CUDA.jl struct ZeRefValue{T} <: Ref{T} diff --git a/test/execution.jl b/test/execution.jl index 0803023e..727c6403 100644 --- a/test/execution.jl +++ b/test/execution.jl @@ -232,6 +232,17 @@ end dims = (16, 16) len = prod(dims) +@testset "CPU array argument" begin + kernel(a) = return + a = oneArray([1f0]) + b = [1f0] + for f in (() -> @oneapi(kernel(b)), () -> a .+ b) + err = try f(); nothing catch err; err end + @test err isa ArgumentError + @test occursin("oneArray", sprint(showerror, err)) + end +end + @testset "manually allocated" begin function kernel(input, output) i = get_global_id()