Hacktoberfest 2026: những issue maintainer đã đánh dấu cho tháng Mười, đang mở và phù hợp người mới. Xem issue Hacktoberfest

Require SPIRVIntrinsics 1.5.1 for local arrays of isbits unions

Đã đóng
  • Có gì không ổn?
#737 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

@maleadt đang làm issue này rồi.

Từ ngày 24/10/2026.

  • #740 của @maleadt — đang mở

Đánh giá

Độ khó
1/5
Thời gian dự kiến
Dưới một giờ
Mức phù hợp với người mới
30/100
Loại issue
Lỗi
Độ rõ ràng
Đặc tả rõ ràng
Mức độ hoạt động
Đình trệ
Lĩnh vực
build-system

Hướng nghiên cứu

Start with the included oneAPI.jl kernel reproducer and run it with @oneapi launch=false to inspect the emitted local-memory allocation. Locate the SPIRVIntrinsics version constraint in the package metadata and require 1.5.1; the kernel should then allocate 68 bytes for the four-element array. Note that linked pull request #740 is already open.

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

Mô tả

SPIRVIntrinsics' emit_localmemory sizes the local memory of a oneLocalArray as len * sizeof(T), plus one selector byte per element for isbits unions. For a union whose size isn't a multiple of its alignment, that's too small: an array of Union{Int64,NTuple{3,Int32}} strides its elements by Base.aligned_sizeof(T) = 16 bytes, not sizeof(T) = 12. JuliaGPU/OpenCL.jl#583 fixes this in SPIRVIntrinsics 1.5.1.

With the current SPIRVIntrinsics 1.5.0, this kernel:

using oneAPI

const T = Union{Int64,NTuple{3,Int32}}

function kernel(out)
    i = get_local_id()
    tmp = oneLocalArray(T, 4)
    tmp[i] = isodd(i) ? Int64(i) : (Int32(i), Int32(i), Int32(i))
    barrier(oneAPI.LOCAL_MEM_FENCE)
    v = tmp[5-i]
    out[i] = v isa Int64 ? v : Int64(v[1]) + 100
    return
end

out = oneArray(zeros(Int64, 4))
@device_code_llvm dump_module=true @oneapi launch=false kernel(out)

allocates 52 bytes of local memory (4 × 12 bytes of data, plus 4 selector bytes):

@local_memory = internal addrspace(3) global [52 x i8] zeroinitializer, align 8

On main, oneDeviceArray also strides by sizeof(T), so the accesses stay within those 52 bytes. But they're misaligned: element 2 is stored at byte offset 12 with an 8-byte alignment assumed:

%15 = mul i64 %2, 12
%16 = getelementptr i8, ptr addrspace(3) @local_memory, i64 %15
store i64 %3, ptr addrspace(3) %16, align 8

#735 makes the stride aligned_sizeof(T), like Array. With it, the elements take 64 bytes and the selector bytes start at byte 64, past the end of the same 52-byte allocation:

%15 = shl i64 %2, 4
%13 = getelementptr i8, ptr addrspace(3) getelementptr (i8, ptr addrspace(3) @local_memory, i64 64), i64 %3

That's an out-of-bounds access of local memory. In this small example the results happen to be correct ([104, 3, 102, 1]), both on main and with #735. With SPIRVIntrinsics 1.5.1, the allocation is 4 * 16 + 4 = 68 bytes. Unions whose size is a multiple of their alignment, like Union{Int,Nothing}, aren't affected.

Ngôn ngữ chính
Julia
Star
218
Fork
36
Merge trung bình
1 ngày 7 giờ
Pull request đã merge (30 ngày)
58

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.