Wrong results from kernels accumulating 2- and 4-byte aggregates on the Aurora LTS stack with Julia 1.13 (IGC 2.11.43)

Start with the reproducer in the issue and the sparse ComplexF16 tests at test/testsuite/sparse.jl:363 and :372; compare behavior across the listed IGC versions and run the reproducer with multiple group and item shapes. The issue identifies IGC as the likely source and lists possible fixes, but no option is tested as a project change. Done means a confirmed fix or workaround makes the failing Julia 1.13 LTS cases correct without breaking relevant tests.

On Aurora's LTS stack with Julia 1.13, a kernel that accumulates a 2- or 4-byte aggregate in a loop silently returns wrong results. This isn't specific to ComplexF16: integer types such as Complex{Int16} and Complex{Int8} fail as well. The first work item of each group is right, and most others are wrong.

This is what keeps test:lts:1.13 red on main. GPUArrays 12.2.0 (JuliaGPU/GPUArrays.jl#796) rewrote the sparse reduction kernels to start from the first stored value, which exposes the bug in the sparse ComplexF16 sum(dx; dims) tests (test/testsuite/sparse.jl:363 and :372). The GPUArrays change only exposes it: a plain @oneapi kernel with the same loop also fails with GPUArrays 12.1.1.

Reproducer

using oneAPI

function accumulate_slices!(out, vals)
    i = get_global_id()
    a = 3i - 2
    v = @inbounds vals[a]
    for k in (a + 1):(a + 2)
        v += @inbounds vals[k]
    end
    @inbounds out[i] = v
    return
end

function check(T; groups, items)
    n = groups * items
    vals = rand(T, 3n)
    out = oneArray{T}(undef, n)
    @oneapi groups = groups items = items accumulate_slices!(out, oneArray(vals))
    oneAPI.synchronize()
    got = Array(out)
    want = [sum(@view vals[(3i - 2):3i]) for i in 1:n]
    return count(i -> !isapprox(got[i], want[i]; atol = 0.01), 1:n)
end

for T in (ComplexF16, Complex{Int16}, Complex{Int8}, ComplexF32)
    println(rpad(T, 16),
        "1 group of 32 items: ", check(T; groups = 1, items = 32), "/32 wrong, ",
        "32 groups of 1 item: ", check(T; groups = 32, items = 1), "/32 wrong")
end

Wrong work items out of 32, with 1 group of 32 items:

With 32 groups of 1 item, all of these are correct, so a test with items = 1 won't catch the bug.

Tested on a Data Center GPU Max 1550 with ONEAPI_LTS=1 and ZE_FLAT_DEVICE_HIERARCHY=FLAT, on NEO 1.6.33578.77 and GPUArrays 12.2.0. For the IGC 2.34.4 runs, only libigc was swapped, and the process's mapped libraries confirmed it was loaded. Only the LTS stack has been tested, not the rolling one.

Which types and kernels are affected

Results are from Julia 1.13 on IGC 2.11.43, using a + loop with shapes 1×32, 4×16 and 32×1.

The shape of the kernel matters as well:

• ComplexF16 is also wrong with a * loop.
• A loop that only replaces the accumulator (if w.a > v.a; v = w; end) is correct for every type above.

These array operations were correct on Julia 1.13 and IGC 2.11.43, for the element types in brackets:

• Dense sum (scalar and dims), and reduce(+, A) (ComplexF16, Float16, Int16, Complex{Int16}).
• extrema (scalar and dims), which holds a pair as its accumulator (Float16, Int16, Int8).
• Sparse CSR and CSC sum(dims) for Float16 and ComplexF32.

The sparse ComplexF16 sum(dims) was the only array operation I tested that gave wrong results. Other operations weren't tested, and any user kernel that accumulates this way can hit it, including KernelAbstractions kernels.

Likely cause

This looks like an IGC bug. On Julia 1.13, the aggregate stays in private memory: the optimized IR for the failing kernels keeps an alloca ([2 x half], i32 or i16), while Julia 1.10 scalarizes it, and nothing is left for IGC to get wrong. Some kernels with a surviving alloca still work ([4 x half], [2 x float]), so the alloca alone doesn't explain it.

• -ze-opt-disable gives correct results.
• A pure SYCL reproducer with a pair of half values in private memory, built with icpx -fsycl -O2, fails the same way on IGC 2.11.43 and is correct on 2.34.4. So Julia isn't required to trigger it.
• In that SYCL reproducer's vISA, the failing kernel has a self-move, mov (M1, 32) V(0,0)<2> V(0,0)<1;1,0>, which matches the mixed-up lanes. This is inferred from the vISA and the wrong values, not from an instruction trace.
• IGC_EnableLdStCombine=0 removes that move and gives correct results in the SYCL reproducer. It hasn't been tested from Julia, and the link to that pass is only a correlation.

I don't know which IGC change between 2.11.43 and 2.34.4 fixed this. It hasn't been reported to Intel yet.

Options

• Get a fixed IGC into the production and LTS stacks. This is the real fix. The full test suite hasn't been run on 2.34.4, only these reproducers and the sparse tests' operations.
• Add an IR cleanup in oneAPI.jl or GPUCompiler that removes the alloca on the LTS stack. This would protect users until IGC is updated, but it's untested.
• For CI only, set IGC_EnableLdStCombine=0 or skip the affected sparse tests on LTS with Julia 1.13. Either one turns main green, but users would still get wrong results.