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

Abierto
  • ¿Algo no está bien?
#757 0 comentarios 0 reacciones 0 asignados Ver en GitHub

Los mantenedores suelen responder en 1 día

Nadie ha tomado este issue todavía.

Evaluación

Dificultad
5/5
Tiempo estimado
Más de una semana
Aptitud para principiantes
35/100
Tipo de issue
Error
Claridad
Bastante claro
Estado de actividad
Activo
Stack tecnológico
julia
Área
backend

Línea de trabajo

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.

Escrito por el modelo de indexación a partir del texto del issue.

Descripción

ci kernels upstream

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:

Element type Julia 1.13, IGC 2.11.43 Julia 1.13, IGC 2.34.4 Julia 1.10, IGC 2.11.43
ComplexF16 31 0 0
Complex{Int16} 31 0 0
Complex{Int8} 31 0 0
ComplexF32 0 0 0

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.

Aggregate size Types tested Result
2 bytes Complex{Int8}, a struct of two Int8 wrong
4 bytes ComplexF16, NTuple{2,Float16}, Complex{Int16}, Complex{UInt16}, and structs of two Float16, two Int16, two UInt16, Float16 + Int16, and four Int8 wrong
6 bytes three Float16 (struct and tuple) correct
8 bytes four Float16 (struct and tuple), four Int16, ComplexF32, Complex{Int32}, and structs of two Float32, two Int32, and Float16 + Float32 correct
16 bytes ComplexF64, and structs of four Float32 and two Float64 correct

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.
Lenguaje dominante
Julia
Estrellas
218
Forks
36
Merge medio
1 d 14 h
PR fusionados (30 d)
59

Preparar el entorno

Este proyecto no incluye contenedor de desarrollo, Dockerfile ni guía de contribución, así que la configuración corre por tu cuenta: empieza por su README y consulta nuestra guía para la primera contribución para los pasos generales.

Primeros pasos

  1. Lee el issue completo y luego la guía de contribución del proyecto.
  2. Comenta en el issue que vas a ocuparte — evita que dos personas hagan lo mismo.
  3. Haz un fork del repositorio y trabaja en una rama.
  4. Abre un pull request que haga referencia al número del issue.

Más de JuliaGPU/oneAPI.jl

Todos los issues de JuliaGPU/oneAPI.jl

Issues similares

Más issues de Julia

Recibe los nuevos issues en tu correo

Un resumen breve de issues de GitHub para principiantes.