diff --git a/Project.toml b/Project.toml index 64be93bf..3f533402 100644 --- a/Project.toml +++ b/Project.toml @@ -59,7 +59,7 @@ PTI_jll = "1.2" PrecompileTools = "1" PrettyTables = "3" Preferences = "1" -SPIRVIntrinsics = "1.4" +SPIRVIntrinsics = "1.5.1" SPIRV_LLVM_Backend_jll = "23" SPIRV_LLVM_Translator_jll = "23" SPIRV_Tools_jll = "2025.4.0" diff --git a/test/unions.jl b/test/unions.jl index 9809a5ff..1da2889c 100644 --- a/test/unions.jl +++ b/test/unions.jl @@ -160,6 +160,31 @@ end @test Array(map(u -> u.x isa Int, d)) == [true, false, true, false] end +@testset "local memory" begin + function kernel_localmemory!(out) + i = get_local_id() + tmp = oneLocalArray(Union{Int64,NTuple{3,Int32}}, 4) + tmp[i] = isodd(i) ? Int64(i) : (Int32(i), Int32(i), Int32(i)) + barrier(oneAPI.LOCAL_MEM_FENCE) + val = tmp[5 - i] + out[i] = val isa Int64 ? val : Int64(val[1]) + 100 + return + end + + out = oneArray(zeros(Int64, 4)) + # The elements have size 12 but stride 16: reserve 64 data bytes and 4 selectors. + # Checking the allocation catches undersizing even if the out-of-bounds accesses + # happen to produce the expected result on a particular driver. + ir = sprint() do io + @device_code_llvm io=io dump_module=true @oneapi launch=false kernel_localmemory!(out) + end + local_sizes = [parse(Int, m[1]) for m in + eachmatch(r"addrspace\(3\) global \[(\d+) x i8\]", ir)] + @test local_sizes == [68] + @oneapi items=4 groups=1 kernel_localmemory!(out) + @test Array(out) == [104, 3, 102, 1] +end + @testset "array operations" begin x = oneArray(Union{Missing,Bool}[missing, true, false]) @test isequal(Array(x .| false), [missing, true, false])