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