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

Đang mở
  • Có gì không ổn?
#757 0 bình luận 0 reaction 0 người được giao Xem trên GitHub

Maintainer thường phản hồi trong vòng 1 ngày

Chưa có ai nhận issue này.

Đánh giá

Độ khó
5/5
Thời gian dự kiến
Hơn một tuần
Mức phù hợp với người mới
35/100
Loại issue
Lỗi
Độ rõ ràng
Khá rõ ràng
Mức độ hoạt động
Sôi nổi
Công nghệ
julia
Lĩnh vực
backend

Hướng nghiên cứu

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.

Do mô hình lập chỉ mục viết ra từ nội dung của issue.

Mô tả

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.
Ngôn ngữ chính
Julia
Star
218
Fork
36
Merge trung bình
1 ngày 14 giờ
Pull request đã merge (30 ngày)
59

Chuẩn bị môi trường

Dự án này không cung cấp dev container, Dockerfile hay hướng dẫn đóng góp, nên bạn cần tự thiết lập môi trường: hãy bắt đầu từ README và xem hướng dẫn đóng góp lần đầu của chúng tôi để biết các bước chung.

Bắt đầu từ đâu

  1. Đọc hết issue, rồi đọc hướng dẫn đóng góp của dự án.
  2. Bình luận trên issue rằng bạn sẽ nhận — tránh hai người làm cùng một việc.
  3. Fork repository và làm thay đổi trên một nhánh.
  4. Mở pull request có tham chiếu số hiệu của issue.

Issue khác của JuliaGPU/oneAPI.jl

Tất cả issue của JuliaGPU/oneAPI.jl

Issue tương tự

Thêm issue về Julia

Nhận issue mới trong hộp thư của bạn

Bản tóm tắt ngắn những issue GitHub phù hợp với người mới.