Correctness issue when indexing "big" oneArrays

I was debugging some Buildkite CI errors in KomaMRI and narrowed the issue down to indexing into a matrix that's “too big.” This is a MWE:

using oneAPI
x = randn(Float32, 513, 3)
x_one = convert(oneArray, x)

julia> y = x[:, 1:2]
513×2 Matrix{Float32}:
 -0.785065    0.605889
 -0.138269   -0.886778
 -0.418909   -0.302653
 -0.0957128  -0.974546
  ⋮          
 -0.486535   -0.594906
  1.42615    -1.00182
  0.0519026  -0.0332484

julia> y_one = x_one[:, 1:2]
513×2 oneArray{Float32, 2, oneAPI.oneL0.DeviceBuffer}:
 -0.785065    0.0
 -0.138269    0.0
 -0.418909    0.0
 -0.0957128   0.0
  ⋮          
 -0.486535    0.0
  1.42615    -1.00182
  0.0519026  -0.0332484

Which is very interesting, as it does not happen when using views:

julia> y = @view x[:, 1:2]
513×2 view(::Matrix{Float32}, :, 1:2) with eltype Float32:
 -0.785065    0.605889
 -0.138269   -0.886778
 -0.418909   -0.302653
 -0.0957128  -0.974546
  ⋮          
 -0.486535   -0.594906
  1.42615    -1.00182
  0.0519026  -0.0332484

julia> y_one = @view x_one[:, 1:2]
513×2 oneArray{Float32, 2, oneAPI.oneL0.DeviceBuffer}:
 -0.785065    0.605889
 -0.138269   -0.886778
 -0.418909   -0.302653
 -0.0957128  -0.974546
  ⋮          
 -0.486535   -0.594906
  1.42615    -1.00182
  0.0519026  -0.0332484

I am using Julia 1.11.6 and oneAPI 2.0.3. Not really sure what could be the cause, but when I use x = randn(Float32, 512, 3) it starts working again, so it seems to be related to the size of the array.

@cncastillo

Tried to reproduce, but it seems to have resolved itself in latest release. Can you confirm @cncastillo?

Using Julia v1.11.7 and oneAPI v2.2.0 I can still can reproduce this on my Linux system (Ubuntu 24.04). I attach the manifest of the environment in case that is useful:

Manifest.toml.txt

EDIT: you might need to increase the size of the array (changing 513 to something larger)

oneAPI v2.2.0 is really old. Can you try with v2.4.0?

Used Julia 1.12.1 and oneAPI v2.4.0 and can still reproduce the error.

The bug was in our KernelAbstractions back-end: @index(Global, Linear) returned the hardware global id, which is wrong when work-groups don't evenly divide a multidimensional range (513 rows → padded 512-item groups), so getindex_kernel skipped elements and wrote out of bounds. It's fixed on main.

We hit this on Aurora (PVC, oneAPI.jl 2.9.2, Julia 1.12.7) in a few more forms, in case it helps anyone searching:

x[:, 2, :, :] on a 4-D oneArray of size (64, 4, 500, B) is wrong for every B ≥ 2, with 768 × (B − 1) device-side Out-of-bounds array access messages. Dropping axis 1 fails the same way; dropping axis 3 or 4 is fine.
rmul!(x, s) on a (64, 1000, 16) ComplexF32 array page-faults and bans the context (@inbounds over a padded ndrange = size(A)), so ifft(x, 2) aborts via AbstractFFTs' ScaledPlan. (256, 1000, 16) is fine because the chosen work-group (256, 4, 1) divides 1000.
permutedims on 4-D arrays is affected the same way.

All of it is the @index(Global, Linear) = get_global_id() issue described above: a CPU re-run of the back-end's launch arithmetic (threads_to_workgroupsize + partition + expand) reproduces every out-of-bounds count exactly.

Is a 2.9.3 with just this fix possible? Our stack (Lux / MLDataDevices / NNlib) still pins GPUArrays 11 and KernelAbstractions 0.9, so we can't move to main until those ecosystems release KernelAbstractions 0.10 support, and in the meantime 2.9.2 silently returns wrong results for common indexing. The change against the v2.9.2 tag is small, mirroring CUDA.jl's back-end:

@device_override @inline function KA.__index_Global_Linear(ctx)
    return _global_linear(ctx, KA.__iterspace(ctx))
end
@inline _global_linear(ctx, ::KA.NDRange{1}) = get_global_id()
@inline function _global_linear(ctx, iterspace::KA.NDRange)
    I = @inbounds KA.expand(iterspace, get_group_id(), get_local_id())
    return @inbounds LinearIndices(KA.__ndrange(ctx))[I]
end

With that patch on top of v2.9.2, all of the cases above are exact on Aurora (PVC, NEO 25.18 LTS), including a direct check that @index(Global, Linear) matches LinearIndices for padded 3-D ranges with automatic and explicit work-group sizes. Happy to open a PR with it plus regression tests (padded N-d ranges with automatic and explicit work-group sizes, and the getindex / permutedims / rmul! cases) against a release-2.9 branch if you'd take one. If a release is not an option, is there a recommended way to pin a fixed oneAPI.jl for KernelAbstractions 0.9 users in the meantime?

@wilkieolin

Thanks @wilkieolin, that analysis is exactly right. Since the fix on main is entangled with the KernelInterface migration (#653, which also moved to KernelAbstractions 0.10 / GPUArrays 12), I cut a release-2.9 branch from v2.9.2 and opened #738 with your patch (as in CUDA.jl, mapping through the iteration space) and regression tests for the cases you listed. Once CI is green I'll merge it and register v2.9.3.

In the meantime you can pin the branch:

] add https://github.com/JuliaGPU/oneAPI.jl#release-2.9

Fix `@index(Global, Linear)` for padded ndranges

Navigation Menu

Repository navigation

Correctness issue when indexing "big" `oneArray`s #509

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Correctness issue when indexing "big" oneArrays #509

Description

Activity

michel2323 commented on Sep 17, 2025

cncastillo commented on Sep 17, 2025

michel2323 commented on Oct 28, 2025

cncastillo commented on Oct 28, 2025

maleadt commented on Oct 19, 2026

wilkieolin commented on Oct 23, 2026

michel2323 commented on Oct 23, 2026

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions

Correctness issue when indexing "big" `oneArray`s #509