Require SPIRVIntrinsics 1.5.1 for local arrays of isbits unions

已关闭
  • 哪里不对？
#737 0 条评论 0 个 reaction 已指派 0 人 在 GitHub 查看

维护者通常 1 天内回复

@maleadt 已经在做这个了。

开始于 2026年10月24日。

  • #740 来自 @maleadt —— 未关闭

评估

难度
1/5
预计耗时
1 小时以内
新手友好度
30/100
Issue 类型
缺陷
描述清晰度
描述清楚
活跃度
停滞
领域
build-system

调研方向

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.

由索引模型根据 Issue 内容生成。

描述

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.

主要语言
Julia
星标
218
派生
36
平均合并
1 天 6 小时
30 天内合并 PR
65

环境准备

这个项目没有提供开发容器、Dockerfile 或贡献指南，环境需要你自己搭建：先看它的 README，通用步骤见我们的新手贡献指南。

从这里开始

  1. 先读完整个 Issue，再读项目的贡献指南。
  2. 在 Issue 下留言说明你要接手 —— 这能避免两个人做同样的事。
  3. Fork 仓库，在一个分支上完成修改。
  4. 提交 Pull Request，并在描述里引用这个 Issue 编号。

JuliaGPU/oneAPI.jl 的其他 Issue

查看 JuliaGPU/oneAPI.jl 的全部 Issue

相似的 Issue

更多 Julia Issue

把新 issue 发到你的邮箱

精选适合新手参与的 GitHub issue 摘要。