Hardware-accelerated linear KA index is wrong

[simeonschaub]

Perhaps I'm misunderstanding something, but I would have expected @index(Global) to always be a valid linear index into indices between 1 and 1856400:

julia> using OpenCL, KernelAbstractions

julia> @kernel function foo(count, indices)
           i = @index(Global)
           k = Atomix.@atomic count[] += 1
           indices[k] = i
       end

julia> count, indices = OpenCL.zeros(Int32), CLArray{Int}(undef, 2000000);

julia> foo(OpenCLBackend())(count, indices; ndrange = (18564, 100))

julia> count
0-dimensional CLArray{Int32, 0, OpenCL.cl.UnifiedDeviceMemory}:
1856400

julia> maximum(indices[OpenCL.@allowscalar 1:count[]])
2046084

[maleadt]

Oh fascinating. This reproduces with a much simpler example:

using OpenCL
import KernelAbstractions as KA

@kernel function foo(max, actual)
    ctx = KA.@context
    I = KA.__index_Global_Cartesian(ctx)
    i = LinearIndices(KA.__ndrange(ctx))[I]
    j = KA.__index_Global_Linear(ctx)

    if j > 20
        @inbounds actual[1] = i
        OpenCL.@println "thread $(get_local_id(1)) block $(get_group_id(1)): I = ($(I[1]), $(I[2])) => i = $i, while j = $j"
    end
end

function main(n = 10, m = 2)
    actual = OpenCL.zeros(Int)
    foo(OpenCLBackend())(n*m, actual; ndrange = (n, m), workgroupsize = 4)
    Array(actual)[1]
end

julia> main()

# some additional output
groups = length(KA.blocks(iterspace)) = 6
items = length(KA.workitems(iterspace)) = 4
global_size = 24
local_size = 4

# kernel output
thread 1 block 6: I = (9, 2) => i = 19, while j = 21
thread 2 block 6: I = (10, 2) => i = 20, while j = 22

So the issue is the global cartesian and linear indices disagreeing. For the (10, 2) ndrange, KA.jl ends up launching 24 threads (with a hard-coded workgroup size of 4). Thread 21 (local id 1, group id 6), which I think should be considered out of bounds, according to KA.jl is pointing to element element (9, 2) with linear index 19. This is because OpenCL does (6-1) * 4 + 1 = 21, while KA.jl does (8-1) * 2 + 1 based on the ndrange of (10, 2). I'm not sure how this is supposed to work, as it's mixing concepts that are incompatible. I'll let @vchuravy chime in when he's back from vacation.

I guess this is also why JuliaGPU/CUDA.jl#2612 fails.