Indexing a StepRange with an Int32 fails to compile to SPIR-V (i128 from the bounds check)

[maleadt]

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.