Hacktoberfest 2026：メンテナが10月に向けて印を付けた、オープンで初心者向けの issue。 Hacktoberfest の issue を見る

Require SPIRVIntrinsics 1.5.1 for local arrays of isbits unions

クローズ
  • 何か問題がありますか？
#737 コメント 0 件 リアクション 0 件 担当者 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時間
マージ済み PR（30日）
65

環境構築

このプロジェクトには開発コンテナ、Dockerfile、コントリビューションガイドがありません。まず README を読み、一般的な手順ははじめてのコントリビューションガイドを参照してください。

はじめの一歩

  1. issue を最後まで読み、次にプロジェクトのコントリビューションガイドを読みます。
  2. 着手することを issue にコメントします — 二人が同じ作業をするのを防げます。
  3. リポジトリをフォークし、ブランチを切って変更します。
  4. issue 番号を参照したプルリクエストを送ります。

JuliaGPU/oneAPI.jl のほかの issue

JuliaGPU/oneAPI.jl の issue をすべて見る

似ている issue

Julia の issue をもっと見る

新しい issue をメールで受け取る

初心者向けの GitHub issue を短くまとめたダイジェスト。