Indexing a StepRange with an Int32 fails to compile to SPIR-V (i128 from the bounds check)
Maintainers usually reply within 1 day
Nobody has claimed this yet.
Assessment
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Newbie friendliness
- 42/100
Research direction
Reproduce the @opencl kernel and the KernelAbstractions copy_kernel! example with --check-bounds=yes. Start with base/range.jl around the StepRange bounds check and GPUCompiler's external_compile path, then inspect the LLVM IR for the i128 arithmetic. Done means the shown Int32-indexed kernels compile to SPIR-V without the arbitrary-precision integer error.
Written by the indexing model from the issue text.
Description
Indexing a StepRange with an Int32 index doesn't compile with the LLVM SPIR-V back-end:
using OpenCL, pocl_jll
function kernel(out, r)
i = get_global_id() % Int32
out[i] = r[i]
return
end
out = CLArray{Int}(undef, 4)
@opencl global_size=4 kernel(out, 1:2:8)
ERROR: LoadError: Failed to compile to SPIR-V with the SPIR-V back-end:
LLVM ERROR: OpTypeInt type with a width other than 8, 16, 32 or 64 bits requires the following SPIR-V extension: SPV_ALTERA_arbitrary_precision_integers
Stacktrace:
[1] error(s::String)
@ Base ./error.jl:44
[2] external_result(...)
@ GPUCompiler ~/.julia/packages/GPUCompiler/sl7Rn/src/utils.jl:120
[3] external_compile(...)
@ GPUCompiler ~/.julia/packages/GPUCompiler/sl7Rn/src/utils.jl:104
...
With i = get_global_id() (an Int), the same kernel works and prints [1, 3, 5, 7].
The i128 comes from Base's bounds check for step ranges, which multiplies in a wider type so that it cannot overflow (base/range.jl):
global function checkbounds(::Type{Bool}, v::StepRange{<:BitInteger64, <:BitInteger64}, i::BitInteger64)
res = widemul(step(v), i-oneunit(i)) + first(v)
(0 < i) & ifelse(0 < step(v), res <= last(v), res >= last(v))
end
The IR handed to the back-end:
%2 = sext i64 %"r::StepRange.step_ptr.unbox" to i128
%3 = sext i32 %0 to i128
%4 = mul nsw i128 %2, %3
%5 = sext i64 %"r::StepRange.unbox" to i128
%6 = add nsw i128 %4, %5
%10 = sext i64 %"r::StepRange.stop_ptr.unbox" to i128
With an Int index, LLVM apparently simplifies the i128 arithmetic away, but not with an Int32 index. KernelAbstractions makes this easy to hit: since JuliaGPU/KernelAbstractions.jl#797, @index is 32-bit on its POCL backend, and a kernel that indexes a view like view(A, 1:2:4, :) fails to compile there under --check-bounds=yes, i.e. in Pkg.test:
using KernelAbstractions # main
@kernel function copy_kernel!(out, v)
i = @index(Global)
out[i] = v[i]
end
A = reshape(collect(Int32, 1:24), 4, 6)
v = view(A, 1:2:4, :)
out = zeros(Int32, length(v))
copy_kernel!(CPU(), 4)(out, v; ndrange=length(v)) # same LLVM ERROR with --check-bounds=yes
AcceleratedKernels' tests run into this for scans and sorts along dims of strided views.
This is probably the same class of problem as #316 (an i2 reaching the SPIR-V module), with a different source of the odd integer width. I haven't checked whether the Khronos translator back-end accepts this.
Versions: Julia 1.13.1, OpenCL.jl 0.10.12, GPUCompiler 2.9.0, LLVM.jl 9.13.2, SPIRV_LLVM_Backend_jll 23.1.1+2, pocl_jll 7.2.0+3 (x86_64 Linux). KernelAbstractions main (d78cad4) shows the same error with pocl_standalone_jll 7.2.0+21.
- Dominant language
- Julia
- Stars
- 313
- Forks
- 50
- Avg merge
- 1d 10h
- Merged PRs (30d)
- 41
Getting set up
This project ships no dev container, Dockerfile or contributing guide, so setting up is up to you: start from its README, and see our first-contribution guide for the general steps.
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
More from JuliaGPU/OpenCL.jl
-
Difficulty 2/5 1-3 hours Newbie friendliness 90/100
Maintainers usually reply within 1 day
-
kernels latency regression
Difficulty 4/5 3-5 days Newbie friendliness 52/100
Maintainers usually reply within 1 day
-
Difficulty 4/5 3-5 days Newbie friendliness 48/100
Maintainers usually reply within 1 day
-
Difficulty 5/5 Over a week Newbie friendliness 35/100
Maintainers usually reply within 1 day
-
Difficulty 5/5 Over a week Newbie friendliness 35/100
Maintainers usually reply within 1 day
All issues in JuliaGPU/OpenCL.jl
Similar issues
-
io
Difficulty 2/5 1-3 hours Newbie friendliness 78/100
Maintainers usually reply within 1 day
-
Difficulty 2/5 1-3 hours Newbie friendliness 72/100
JuliaGPU/GPUArrays.jl#837 ·
Maintainers usually reply within 2 days
-
Difficulty 1/5 Under an hour Newbie friendliness 88/100
Maintainers usually reply within 1 day
-
Difficulty 1/5 Under an hour Newbie friendliness 68/100
-
Difficulty 2/5 1-3 hours Newbie friendliness 78/100
bmad-sim/BeamTracking.jl#382 ·
Maintainers usually reply within 1 day