From 262e6c980a72bfc0c8dc7b1762a72d658df68cae Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Fri, 23 Oct 2026 15:03:33 +0200 Subject: [PATCH] Support arrays of isbits unions oneArray rejected element types like `Union{Int,Nothing}` or `Union{Missing,Bool}`, although the device side already knew how to load and store them. Store them like `Array` does: the data, padded to `aligned_sizeof` per element, followed by one selector byte per element. `maxsize` locates an array's selector bytes. Derived arrays advance it by their element offset, which the byte offset can't recover for unions of singletons, so kernels on views of these arrays use the view's own selector bytes. Every copy path (to and from the host, between device arrays, and through host pointers into shared and host memory) and `resize!` copy the selector bytes too; scalar indexing of host-accessible arrays goes through such a copy, and `unsafe_wrap(Array, ...)` rejects unions as an `Array` keeps its selector bytes elsewhere. On the device, `Const` arrays (used by AcceleratedKernels' reductions) load unions through their selector byte. Structs with a union field took the union path too, reading a selector byte past the end of the array; they keep their selector inside the element, and now load as plain bits. --- src/array.jl | 143 +++++++++++++++++------------ src/device/array.jl | 27 +++--- src/memory.jl | 2 +- src/oneAPI.jl | 9 ++ test/unions.jl | 214 ++++++++++++++++++++++++++++++++++++++++++++ 5 files changed, 329 insertions(+), 66 deletions(-) create mode 100644 test/unions.jl diff --git a/src/array.jl b/src/array.jl index 3f968c2a5..f5b063815 100644 --- a/src/array.jl +++ b/src/array.jl @@ -65,7 +65,6 @@ end function check_eltype(T) Base.allocatedinline(T) || error("oneArray only supports element types that are stored inline") - Base.isbitsunion(T) && error("oneArray does not yet support isbits-union arrays") support = eltype_support() if !support.float16 contains_eltype(T, Float16) && error("Float16 is not supported on this device") @@ -87,7 +86,7 @@ end N-dimensional dense array type for Intel GPU programming using oneAPI and Level Zero. # Type Parameters -- `T`: Element type (must be stored inline, no isbits-unions) +- `T`: Element type (must be stored inline) - `N`: Number of dimensions - `B`: Buffer type, one of: - `oneL0.DeviceBuffer`: GPU device memory (default, not CPU-accessible) @@ -124,23 +123,21 @@ See also: [`oneVector`](@ref), [`oneMatrix`](@ref), [`is_device`](@ref), [`is_sh mutable struct oneArray{T,N,B} <: AbstractGPUArray{T,N} data::DataRef{Managed{B}} - maxsize::Int # maximum data size; excluding any selector bytes + # size of the data in the buffer, in bytes. arrays of isbits unions keep one selector byte + # per element past the data; for those, this is where the array's first selector byte + # is (see `typetagdata`). + maxsize::Int offset::Int # offset of the data in the buffer, in bytes dims::Dims{N} function oneArray{T,N,B}(::UndefInitializer, dims::Dims{N}) where {T,N,B} check_eltype(T) - maxsize = prod(dims) * sizeof(T) - bufsize = if Base.isbitsunion(T) - # type tag array past the data - maxsize + prod(dims) - else - maxsize - end + maxsize = prod(dims) * aligned_sizeof(T) + bufsize = Base.isbitsunion(T) ? maxsize + prod(dims) : maxsize ctx = context() dev = device() - alignment = Base.datatype_alignment(T) + alignment = element_alignment(T) data = GPUArrays.cached_alloc((oneArray, B, ctx, dev, bufsize, alignment)) do DataRef(release, alloc_managed(B, ctx, dev, bufsize, alignment)) end @@ -149,11 +146,13 @@ mutable struct oneArray{T,N,B} <: AbstractGPUArray{T,N} end function oneArray{T,N}(data::DataRef{Managed{B}}, dims::Dims{N}; - maxsize::Int=prod(dims) * sizeof(T), offset::Int=0) where {T,N,B} + maxsize::Int=prod(dims) * aligned_sizeof(T), + offset::Int=0) where {T,N,B} check_eltype(T) if sizeof(T) == 0 offset == 0 || error("Singleton arrays cannot have a nonzero offset") - maxsize == 0 || error("Singleton arrays cannot have a size") + # (a union of singletons still needs `maxsize` to find its selector bytes) + Base.isbitsunion(T) || maxsize == 0 || error("Singleton arrays cannot have a size") end obj = new{T,N,B}(copy(data), maxsize, offset, dims) finalizer(unsafe_free!, obj) @@ -162,6 +161,10 @@ end GPUArrays.storage(a::oneArray) = a.data +# the alignment of the elements of an array of `T` +element_alignment(::Type{T}) where {T} = + Base.isbitsunion(T) ? Base.uniontype_layout(T)[3] : Base.datatype_alignment(T) + ## alias detection @@ -235,7 +238,7 @@ end ## array interface -Base.elsize(::Type{<:oneArray{T}}) where {T} = sizeof(T) +Base.elsize(::Type{<:oneArray{T}}) where {T} = aligned_sizeof(T) Base.size(x::oneArray) = x.dims Base.sizeof(x::oneArray) = Base.elsize(x) * length(x) @@ -392,13 +395,23 @@ device_pointer(x::oneArray{T}, i::Integer=1) where {T} = # Host-accessible arrays can be indexed from CPU, bypassing GPUArrays restrictions. # Taking the host pointer waits for outstanding work on the memory. -function Base.getindex(x::oneArray{<:Any, <:Any, <:Union{oneL0.HostBuffer, oneL0.SharedBuffer}}, I::Int) +function Base.getindex(x::oneArray{T, <:Any, <:Union{oneL0.HostBuffer, oneL0.SharedBuffer}}, I::Int) where {T} @boundscheck checkbounds(x, I) + if Base.isbitsunion(T) + # a union can't be loaded from a pointer; copy the element and its selector byte + y = Array{T}(undef, 1) + copyto!(y, 1, x, I, 1) + return y[1] + end return unsafe_load(pointer(x, I; type = oneL0.HostBuffer)) end -function Base.setindex!(x::oneArray{<:Any, <:Any, <:Union{oneL0.HostBuffer, oneL0.SharedBuffer}}, v, I::Int) +function Base.setindex!(x::oneArray{T, <:Any, <:Union{oneL0.HostBuffer, oneL0.SharedBuffer}}, v, I::Int) where {T} @boundscheck checkbounds(x, I) + if Base.isbitsunion(T) + copyto!(x, I, T[v], 1, 1) + return x + end return unsafe_store!(pointer(x, I; type = oneL0.HostBuffer), v) end @@ -413,12 +426,30 @@ end ## memory copying +# arrays of isbits unions keep one selector byte per element past the data, like `Array` +# does: the zero-based index of the element's type in `Base.uniontypes(T)`. + +@static if VERSION >= v"1.11.0-DEV.753" +function typetagdata(a::Array, i=1) + # for memory of isbits unions, the reference holds the offset of the data in elements + ptr_or_offset = Int(a.ref.ptr_or_offset) + @ccall(jl_genericmemory_typetagdata(a.ref.mem::Any)::Ptr{UInt8}) + ptr_or_offset + i - 1 +end +else typetagdata(a::Array, i=1) = ccall(:jl_array_typetagdata, Ptr{UInt8}, (Any,), a) + i - 1 -function typetagdata(a::oneArray, i=1) - # for zero-size element types (e.g. singleton unions), the byte offset - # is always zero, so the corresponding element offset is also zero - elem_offset = iszero(Base.elsize(a)) ? 0 : a.offset ÷ Base.elsize(a) - return convert(ZePtr{UInt8}, a.data[].mem) + a.maxsize + elem_offset + i - 1 +end + +# the selector byte of element `i`. the host pointer of a host-accessible array doesn't wait +# for outstanding work, unlike `pointer`. +function typetagdata(a::oneArray, i=1; type=oneL0.DeviceBuffer) + PT = if type == oneL0.DeviceBuffer + ZePtr{UInt8} + elseif type == oneL0.HostBuffer + Ptr{UInt8} + else + error("unknown memory type") + end + return convert(PT, a.data[].mem) + a.maxsize + i - 1 end function Base.copyto!(dest::oneArray{T}, doffs::Integer, src::Array{T}, soffs::Integer, @@ -467,12 +498,13 @@ Base.copyto!(dest::oneDenseArray{T}, src::oneDenseArray{T}) where {T} = function Base.unsafe_copyto!(ctx::ZeContext, dev::ZeDevice, dest::oneDenseArray{T}, doffs, src::Array{T}, soffs, n) where T stream = global_stream(ctx, dev) - bytes = n * sizeof(T) + bytes = n * aligned_sizeof(T) GC.@preserve src dest begin # the copy reads the source when it executes, by which time the source may have # changed or been freed. small uploads copy it first if the device stages them and - # there's room, others wait for the copy. - staged = 0 < bytes <= MAX_STAGED && stages_uploads(dev) && + # there's room, others wait for the copy. (unions aren't staged, as their selector + # bytes would need a staged copy of their own.) + staged = 0 < bytes <= MAX_STAGED && !Base.isbitsunion(T) && stages_uploads(dev) && with_ownership(stream, Managed[dest.data[]]) do execute!(stream) do list stage_upload!(list, ctx, dev, device_pointer(dest, doffs), pointer(src, soffs), bytes) @@ -481,14 +513,13 @@ function Base.unsafe_copyto!(ctx::ZeContext, dev::ZeDevice, if !staged with_ownership(stream, Managed[dest.data[]]) do unsafe_copyto!(ctx, dev, device_pointer(dest, doffs), pointer(src, soffs), n) + if Base.isbitsunion(T) + unsafe_copyto!(ctx, dev, typetagdata(dest, doffs), typetagdata(src, soffs), n) + end end synchronize(stream) end end - if Base.isbitsunion(T) - # copy selector bytes - error("oneArray does not yet support isbits-union arrays") - end return dest end @@ -499,12 +530,11 @@ function Base.unsafe_copyto!(ctx::ZeContext, dev::ZeDevice, GC.@preserve src dest begin with_ownership(stream, Managed[src.data[]]) do unsafe_copyto!(ctx, dev, pointer(dest, doffs), device_pointer(src, soffs), n) + if Base.isbitsunion(T) + unsafe_copyto!(ctx, dev, typetagdata(dest, doffs), typetagdata(src, soffs), n) + end end end - if Base.isbitsunion(T) - # copy selector bytes - error("oneArray does not yet support isbits-union arrays") - end # copies to the host are synchronizing synchronize(stream) @@ -517,12 +547,11 @@ function Base.unsafe_copyto!(ctx::ZeContext, dev::ZeDevice, GC.@preserve src dest begin with_ownership(global_stream(ctx, dev), Managed[dest.data[], src.data[]]) do unsafe_copyto!(ctx, dev, device_pointer(dest, doffs), device_pointer(src, soffs), n) + if Base.isbitsunion(T) + unsafe_copyto!(ctx, dev, typetagdata(dest, doffs), typetagdata(src, soffs), n) + end end end - if Base.isbitsunion(T) - # copy selector bytes - error("oneArray does not yet support isbits-union arrays") - end return dest end @@ -530,14 +559,13 @@ end function Base.unsafe_copyto!(ctx::ZeContext, dev::ZeDevice, dest::oneDenseArray{T,<:Any,<:Union{oneL0.SharedBuffer,oneL0.HostBuffer}}, doffs, src::Array{T}, soffs, n) where T - if Base.isbitsunion(T) - # copy selector bytes - error("oneArray does not yet support isbits-union arrays") - end # the host pointer waits for outstanding work; hold the lock so that no new work is # submitted while copying GC.@preserve src dest with_managed_locks(Managed[dest.data[]]) do unsafe_copyto!(pointer(dest, doffs; type=oneL0.HostBuffer), pointer(src, soffs), n) + if Base.isbitsunion(T) + unsafe_copyto!(typetagdata(dest, doffs; type=oneL0.HostBuffer), typetagdata(src, soffs), n) + end end return dest @@ -545,14 +573,13 @@ end function Base.unsafe_copyto!(ctx::ZeContext, dev::ZeDevice, dest::Array{T}, doffs, src::oneDenseArray{T,<:Any,<:Union{oneL0.SharedBuffer,oneL0.HostBuffer}}, soffs, n) where T - if Base.isbitsunion(T) - # copy selector bytes - error("oneArray does not yet support isbits-union arrays") - end # the host pointer waits for outstanding work; hold the lock so that no new work is # submitted while copying GC.@preserve src dest with_managed_locks(Managed[src.data[]]) do unsafe_copyto!(pointer(dest, doffs), pointer(src, soffs; type=oneL0.HostBuffer), n) + if Base.isbitsunion(T) + unsafe_copyto!(typetagdata(dest, doffs), typetagdata(src, soffs; type=oneL0.HostBuffer), n) + end end return dest @@ -605,11 +632,15 @@ function GPUArrays.derive(::Type{T}, a::oneArray, dims::Dims{N}, offset::Int) wh if sizeof(T) == 0 Base.elsize(a) == 0 || error("Cannot derive a singleton array from non-singleton inputs") end - offset = a.offset + offset * sizeof(T) + # the selector bytes are indexed by element, so advance them here: the byte offset can't + # tell how many elements of zero size (like unions of singletons) were skipped. arrays of + # unions only derive arrays of the same element type, as `reinterpret` rejects them. + maxsize = Base.isbitsunion(T) ? a.maxsize + offset : a.maxsize + offset = a.offset + offset * aligned_sizeof(T) # The derived array constructor copies `a.data`, but merely loading that field does not # keep `a` alive. Without this preserve, `a` may be finalized between the field load and # the DataRef copy, causing its finalizer to mark the reference as freed. - GC.@preserve a oneArray{T,N}(a.data, dims; a.maxsize, offset) + GC.@preserve a oneArray{T,N}(a.data, dims; maxsize, offset) end @@ -625,7 +656,7 @@ function Base.unsafe_convert(::Type{ZePtr{T}}, V::SubArray{T,N,P,<:Tuple{Vararg{ end function Base.unsafe_convert(::Type{ZePtr{T}}, V::SubArray{T,N,P,<:Tuple{Vararg{Union{Base.RangeIndex,Base.ReshapedUnitRange}}}}) where {T,N,P} return Base.unsafe_convert(ZePtr{T}, parent(V)) + - (Base.first_index(V)-1)*sizeof(T) + (Base.first_index(V)-1)*aligned_sizeof(T) end @@ -659,6 +690,9 @@ a shared buffer (as created by `oneArray{T}(undef, ...)`) or a host buffer. """ function Base.unsafe_wrap(::Type{Array}, arr::oneArray{T,N,<:Union{oneL0.SharedBuffer,oneL0.HostBuffer}}) where {T,N} + # an `Array` keeps its selector bytes elsewhere than a `oneArray` does + Base.isbitsunion(T) && throw(ArgumentError( + "unsafe_wrap(Array, ...) does not support isbits-union element types like $T")) # TODO: can we make this more convenient by increasing the buffer's refcount and using # a finalizer on the Array? does that work when taking views etc of the Array? ptr = pointer(arr; type=oneL0.HostBuffer) @@ -677,24 +711,23 @@ guaranteed to be initialized. """ function Base.resize!(a::oneVector{T}, n::Integer) where {T} # TODO: add additional space to allow for quicker resizing - maxsize = n * sizeof(T) - bufsize = if isbitstype(T) - maxsize - else - # type tag array past the data - maxsize + n - end + maxsize = n * aligned_sizeof(T) + bufsize = Base.isbitsunion(T) ? maxsize + n : maxsize # replace the data with a new one. this 'unshares' the array. # as a result, we can safely support resizing unowned buffers. ctx = context(a) dev = device(a) new_data = DataRef(release, alloc_managed(buftype(a), ctx, dev, bufsize, - Base.datatype_alignment(T))) + element_alignment(T))) m = min(length(a), n) if m > 0 with_ownership(global_stream(ctx, dev), Managed[new_data[], a.data[]]) do unsafe_copyto!(ctx, dev, convert(ZePtr{T}, new_data[].mem), device_pointer(a), m) + if Base.isbitsunion(T) + unsafe_copyto!(ctx, dev, convert(ZePtr{UInt8}, new_data[].mem) + maxsize, + typetagdata(a), m) + end end end unsafe_free!(a) diff --git a/src/device/array.jl b/src/device/array.jl index 75ba79088..b948da7ec 100644 --- a/src/device/array.jl +++ b/src/device/array.jl @@ -44,6 +44,8 @@ See also: [`oneArray`](@ref), [`oneLocalArray`](@ref), [`@oneapi`](@ref) """ struct oneDeviceArray{T,N,A} <: DenseArray{T,N} ptr::LLVMPtr{T,A} + # distance from `ptr` to the selector bytes of an array of isbits unions (see + # `typetagdata`); the size of the data for an array that starts at its memory maxsize::Int dims::Dims{N} @@ -52,7 +54,7 @@ struct oneDeviceArray{T,N,A} <: DenseArray{T,N} # inner constructors, fully parameterized, exact types (ie. Int not <:Integer) # TODO: deprecate; put `ptr` first like oneArray oneDeviceArray{T,N,A}(dims::Dims{N}, ptr::LLVMPtr{T,A}, - maxsize::Int=prod(dims)*sizeof(T)) where {T,A,N} = + maxsize::Int=prod(dims)*aligned_sizeof(T)) where {T,A,N} = new(ptr, maxsize, dims, prod(dims)) end @@ -76,7 +78,7 @@ oneDeviceVector{T,A}(len::Integer, p::LLVMPtr{T,A}) where {T,A} ## array interface -Base.elsize(::Type{<:oneDeviceArray{T}}) where {T} = sizeof(T) +Base.elsize(::Type{<:oneDeviceArray{T}}) where {T} = aligned_sizeof(T) Base.size(g::oneDeviceArray) = g.dims Base.sizeof(x::oneDeviceArray) = Base.elsize(x) * length(x) @@ -90,6 +92,7 @@ Base.pointer(x::oneDeviceArray{T,<:Any,A}) where {T,A} = Base.unsafe_convert(LLV Base.unsafe_convert(LLVMPtr{T,A}, x) + Base._memory_offset(x, i) end +# the selector byte of element `i` of an array of isbits unions typetagdata(a::oneDeviceArray{<:Any,<:Any,A}, i=1) where {A} = reinterpret(LLVMPtr{UInt8,A}, a.ptr + a.maxsize) + i - one(i) @@ -121,10 +124,10 @@ end #@boundscheck checkbounds(A, index) @boundscheck index <= length(A) || Base.throw_boundserror(A, index) - if isbitstype(T) - arrayref_bits(A, index) - else #if isbitsunion(T) + if Base.isbitsunion(T) arrayref_union(A, index) + else + arrayref_bits(A, index) end end @@ -166,10 +169,10 @@ end #@boundscheck checkbounds(A, index) @boundscheck index <= length(A) || Base.throw_boundserror(A, index) - if isbitstype(T) - arrayset_bits(A, x, index) - else #if isbitsunion(T) + if Base.isbitsunion(T) arrayset_union(A, x, index) + else + arrayset_bits(A, x, index) end return A end @@ -207,8 +210,12 @@ end #@boundscheck checkbounds(A, index) @boundscheck index <= length(A) || Base.throw_boundserror(A, index) - align = alignment(A) - unsafe_cached_load(pointer(A), index, Val(align)) + if Base.isbitsunion(T) + arrayref_union(A, index) + else + align = alignment(A) + unsafe_cached_load(pointer(A), index, Val(align)) + end end diff --git a/src/memory.jl b/src/memory.jl index 4afeabdf2..a5f28f545 100644 --- a/src/memory.jl +++ b/src/memory.jl @@ -23,7 +23,7 @@ See also: [`copyto!`](@ref), [`oneArray`](@ref) """ function Base.unsafe_copyto!(ctx::ZeContext, dev::ZeDevice, dst::Union{Ptr{T},ZePtr{T}}, src::Union{Ptr{T},ZePtr{T}}, N::Integer) where T - bytes = N*sizeof(T) + bytes = N*aligned_sizeof(T) bytes==0 && return execute!(global_stream(ctx, dev)) do list append_copy!(list, dst, src, bytes) diff --git a/src/oneAPI.jl b/src/oneAPI.jl index a55a550f4..941b17045 100644 --- a/src/oneAPI.jl +++ b/src/oneAPI.jl @@ -36,6 +36,15 @@ include("../lib/level-zero/oneL0.jl") using .oneL0 functional() = oneL0.functional[] +# the size of an array element. this differs from `sizeof` for isbits unions whose size isn't +# a multiple of their alignment, e.g., `Union{Int64,NTuple{3,Int32}}` (12 vs. 16 bytes). +@static if VERSION >= v"1.11" + import Base: aligned_sizeof +else + # Julia 1.10's doesn't fold to a constant in device code + @generated aligned_sizeof(::Type{T}) where {T} = :($(Base.aligned_sizeof(T))) +end + # device functionality import SPIRVIntrinsics SPIRVIntrinsics.@import_all diff --git a/test/unions.jl b/test/unions.jl new file mode 100644 index 000000000..9809a5ff6 --- /dev/null +++ b/test/unions.jl @@ -0,0 +1,214 @@ +import KernelInterface as KI + +# Arrays of isbits unions keep one selector byte per element past the data (see +# `typetagdata` in src/array.jl), like `Array` does. + +const buffer_types = (oneL0.DeviceBuffer, oneL0.SharedBuffer, oneL0.HostBuffer) + +# a struct with a union field is stored inline, with its selector inside the element +struct UnionField + x::Union{Int,Float32} +end + +# compare the elements' types too, which `isequal` ignores (e.g. `isequal(Int8(1), 1)`) +same(x, y) = size(x) == size(y) && all(map(===, x, y)) + +# increment the integer elements of `x` +function bump!(x) + i = get_global_id() + if i <= length(x) + val = x[i] + if val isa Int + x[i] = val + 1 + end + end + return +end + +@testset "size" begin + a = oneArray(Union{Int,Nothing}[1, nothing, 3]) + @test sizeof(a) == 3 * sizeof(Int) # excludes the selector bytes + + # members smaller than the union's alignment get padded + T = Union{Int64,NTuple{3,Int32}} + b = oneArray(T[1, (Int32(2), Int32(3), Int32(4))]) + @test sizeof(b) == 2 * 16 + @test pointer(b, 2) - pointer(b, 1) == 16 + @test pointer(view(b, reshape(2:2, 1, 1))) == pointer(b, 2) +end + +@testset "round trips ($B)" for B in buffer_types + for a in (Union{Int,Nothing}[1, nothing, 3], + Union{Int8,Int64}[1, Int8(2), 3], + Union{NTuple{3,Int8},Int16}[Int8.((1, 2, 3)), Int16(4), Int8.((5, 6, 7))], + Union{Int64,NTuple{3,Int32}}[1, (Int32(2), Int32(3), Int32(4)), 5], + Union{Nothing,Missing}[nothing, missing, missing, nothing], + Union{Union{Missing,Nothing},Int}[0, missing, nothing]) + T = eltype(a) + b = oneArray{T,1,B}(a) + @test same(Array(b), a) + @test same(Array(copy(b)), a) + @test same(Array(similar(b) .= b), a) + + # between device arrays, also of other buffer types + for B′ in buffer_types + c = oneArray{T,1,B′}(undef, length(a)) + copyto!(c, b) + @test same(Array(c), a) + copyto!(c, 2, b, 1, 2) + @test same(Array(c)[1:3], T[a[1], a[1], a[2]]) + end + + # from and to offsets in host arrays + d = oneArray{T,1,B}(undef, 2) + copyto!(d, 1, a, 2, 2) + @test same(Array(d), a[2:3]) + e = similar(a) + copyto!(e, 1, a, 1, 1) + copyto!(e, 2, d, 1, 2) + @test same(e[1:3], a[1:3]) + end +end + +@testset "scalar indexing ($B)" for B in buffer_types + a = oneArray{Union{Int,Nothing},1,B}([1, nothing, 3]) + @allowscalar begin + @test a[1] === 1 + @test a[2] === nothing + a[1] = nothing + a[2] = 2 + end + @test Array(a) == [nothing, 2, 3] +end + +@testset "views" begin + # the parent's first selector bytes differ from the view's + a = Union{Int,Nothing}[nothing, nothing, 1, nothing, 3, 5] + b = oneArray(a) + v = view(b, 3:5) + @test Array(v) == a[3:5] + @test Array(copy(v)) == a[3:5] + @test @allowscalar v[2] === nothing + + # kernels see the view's elements, and leave the parent alone outside of it + @oneapi items=length(v) bump!(v) + @test Array(v) == [2, nothing, 4] + @test Array(b) == [nothing, nothing, 2, nothing, 4, 5] + + # views of views, and reshaped views + w = view(v, 2:3) + @test Array(w) == [nothing, 4] + @oneapi items=length(w) bump!(w) + @test Array(b) == [nothing, nothing, 2, nothing, 5, 5] + r = reshape(view(b, 3:6), 2, 2) + @test Array(r) == [2 5; nothing 5] + @oneapi items=length(r) bump!(r) + @test Array(b) == [nothing, nothing, 3, nothing, 6, 6] + + # copies into a view + copyto!(view(b, 1:2), Union{Int,Nothing}[7, 8]) + @test Array(b) == [7, 8, 3, nothing, 6, 6] + copyto!(view(b, 3:4), view(b, 1:2)) + @test Array(b) == [7, 8, 7, 8, 6, 6] + view(b, 5:6) .= nothing + @test Array(b) == [7, 8, 7, 8, nothing, nothing] + + # non-contiguous views index into their parent + nc = view(b, 1:2:5) + @oneapi items=length(nc) bump!(nc) + @test Array(b) == [8, 8, 8, 8, nothing, nothing] + + # unions of singletons have no data, so their views don't have a byte offset either + c = oneArray([nothing, nothing, missing, missing]) + @test isequal(Array(view(c, 3:4)), [missing, missing]) + @test isequal(Array(copy(view(c, 2:3))), [nothing, missing]) + function ismissing_kernel(ok, x) + i = get_global_id() + if i <= length(x) + ok[i] = x[i] === missing + end + return + end + ok = oneArray(fill(false, 2)) + @oneapi items=2 ismissing_kernel(ok, view(c, 3:4)) + @test all(Array(ok)) +end + +@testset "kernels" begin + a = oneArray(Union{Int,Nothing}[1, nothing, 3]) + @oneapi items=length(a) bump!(a) + @test Array(a) == [2, nothing, 4] + + b = oneArray(Union{Union{Missing,Nothing},Int}[0, missing, nothing]) + @oneapi items=length(b) bump!(b) + @test isequal(Array(b), [1, missing, nothing]) + + c = oneArray(Union{Int64,NTuple{3,Int32}}[1, (Int32(2), Int32(3), Int32(4)), 5]) + @oneapi items=length(c) bump!(c) + @test Array(c) == Union{Int64,NTuple{3,Int32}}[2, (Int32(2), Int32(3), Int32(4)), 6] + + function fill_struct!(x) + i = get_global_id() + if i <= length(x) + x[i] = UnionField(isodd(i) ? Int(i) : Float32(i)) + end + return + end + d = oneArray{UnionField}(undef, 4) + @oneapi items=4 fill_struct!(d) + @test Array(d) == [UnionField(1), UnionField(2f0), UnionField(3), UnionField(4f0)] + @test Array(map(u -> u.x isa Int, d)) == [true, false, true, false] +end + +@testset "array operations" begin + x = oneArray(Union{Missing,Bool}[missing, true, false]) + @test isequal(Array(x .| false), [missing, true, false]) + @test isequal(Array(map(!, x)), [missing, false, true]) + @test any(x) === true + @test all(x) === false + @test isequal(Array(reverse(x)), [false, true, missing]) + @test isequal(Array(vcat(x, x)), vcat(Array(x), Array(x))) + + y = similar(x) + fill!(y, missing) + @test all(ismissing, Array(y)) + fill!(y, true) + @test Array(y) == [true, true, true] + + z = oneArray(Union{Int,Nothing}[1, nothing, 3, nothing]) + @test sum(something.(z, 0)) == 4 + @test mapreduce(v -> something(v, 10), +, z) == 24 + @test Array(reshape(z, 2, 2)) == [1 3; nothing nothing] + + w = similar(z) + KI.copyto!(oneAPIBackend(), w, z) + @test Array(w) == Array(z) +end + +@testset "resize!" begin + a = oneArray(Union{Int,Nothing}[1, nothing, 3]) + resize!(a, 5) + @test length(a) == 5 + @test Array(a)[1:3] == [1, nothing, 3] + resize!(a, 2) + @test Array(a) == [1, nothing] + + # arrays derived from a view start at an offset + b = vec(view(oneArray(Union{Int,Nothing}[nothing, 1, nothing, 3]), 2:4)) + resize!(b, 4) + @test Array(b)[1:3] == [1, nothing, 3] + + c = oneArray([nothing, missing, missing]) + resize!(c, 4) + @test isequal(Array(c)[1:3], [nothing, missing, missing]) + resize!(c, 1) + @test Array(c) == [nothing] + d = vec(view(oneArray([nothing, missing, missing]), 2:3)) + resize!(d, 1) + @test isequal(Array(d), [missing]) +end + +@testset "unsupported operations" begin + a = oneArray{Union{Int,Nothing},1,oneL0.SharedBuffer}([1, nothing]) + @test_throws ArgumentError unsafe_wrap(Array, a) +end