From ce7bd83cb5a5c0569e8fc1a2ef5b669e109b2ddd Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Tue, 29 Sep 2026 21:02:34 +0200 Subject: [PATCH] PoCL: compute the typed index queries with `% T` `T(x)` is a checked conversion, which leaves a `throw_inexacterror` branch in every kernel using a typed index query. `% T` doesn't, but needs SPIRVIntrinsics 1.1.3 (JuliaGPU/OpenCL.jl#511): with older versions, InstCombine folds the truncated 3-D builtin load into a vector type SPIR-V can't represent. --- Project.toml | 2 +- src/pocl/backend.jl | 15 +++++++++------ 2 files changed, 10 insertions(+), 7 deletions(-) diff --git a/Project.toml b/Project.toml index 48c72ffb8..4d02e6782 100644 --- a/Project.toml +++ b/Project.toml @@ -51,7 +51,7 @@ Printf = "<0.0.1, 1" Random = "1" Random123 = "1.7.1" RandomNumbers = "1.6.0" -SPIRVIntrinsics = "1" +SPIRVIntrinsics = "1.1.3" SPIRV_LLVM_Backend_jll = "23" SPIRV_Tools_jll = "2024.4, 2025.1" SparseArrays = "<0.0.1, 1.6" diff --git a/src/pocl/backend.jl b/src/pocl/backend.jl index fc0d4b332..9e9c4177f 100644 --- a/src/pocl/backend.jl +++ b/src/pocl/backend.jl @@ -282,28 +282,31 @@ end ## Indexing Functions +# `% T` rather than `T(x)`: a checked conversion leaves an error branch in every kernel. +# This needs SPIRVIntrinsics 1.1.3, whose 3-D builtins survive the truncation. + @device_override @inline function KI.get_local_id(::Type{T}) where {T} - return (; x = T(get_local_id(1)), y = T(get_local_id(2)), z = T(get_local_id(3))) + return (; x = get_local_id(1) % T, y = get_local_id(2) % T, z = get_local_id(3) % T) end @device_override @inline function KI.get_group_id(::Type{T}) where {T} - return (; x = T(get_group_id(1)), y = T(get_group_id(2)), z = T(get_group_id(3))) + return (; x = get_group_id(1) % T, y = get_group_id(2) % T, z = get_group_id(3) % T) end @device_override @inline function KI.get_global_id(::Type{T}) where {T} - return (; x = T(get_global_id(1)), y = T(get_global_id(2)), z = T(get_global_id(3))) + return (; x = get_global_id(1) % T, y = get_global_id(2) % T, z = get_global_id(3) % T) end @device_override @inline function KI.get_local_size(::Type{T}) where {T} - return (; x = T(get_local_size(1)), y = T(get_local_size(2)), z = T(get_local_size(3))) + return (; x = get_local_size(1) % T, y = get_local_size(2) % T, z = get_local_size(3) % T) end @device_override @inline function KI.get_num_groups(::Type{T}) where {T} - return (; x = T(get_num_groups(1)), y = T(get_num_groups(2)), z = T(get_num_groups(3))) + return (; x = get_num_groups(1) % T, y = get_num_groups(2) % T, z = get_num_groups(3) % T) end @device_override @inline function KI.get_global_size(::Type{T}) where {T} - return (; x = T(get_global_size(1)), y = T(get_global_size(2)), z = T(get_global_size(3))) + return (; x = get_global_size(1) % T, y = get_global_size(2) % T, z = get_global_size(3) % T) end @device_override KI.get_sub_group_size() = get_sub_group_size() % UInt32