`ndrange` provided in KernelAbstractions kernels is broken

[VarLad]

I'm unable to pinpoint at what point the error is occurring, as debugging KernelAbstractions kernels is different from normal Julia, so some guidance here would be helpful.
The issue is that some elements of the output array are 0.

I can't test this with any other driver other than PoCL either, since NVIDIA doesn't support SPIRV, so it would be nice if someone could test if it gives a different behavior with Intel drivers.
This probably points to a problem in our current code somewhere. Behavior is same both before and after the USM PR.

Edit:

Most recently:

Here's a reproducer:

using OpenCL, pocl_jll, KernelAbstractions

@kernel inbounds=true function _mwe!(@Const(v))
           temp = @localmem Int8 (1,)
           i = @index(Global, Linear)
           @print i "\n"
           @synchronize()
       end

v = CLArray(rand(Float32, 10))

_mwe!(OpenCLBackend(), 256)(v, ndrange=length(v))

This prints 1...256.

The CUDA version of the same code prints 1...10.

using CUDA

b = CuArray(rand(Float32, 10))

_mwe!(CUDABackend(false, false), 256)(b, ndrange=length(b))

The issue is probably that ndrange is not working.

Consequently, creating a CLArray of size (multiple of) 256 works without any issues, for the any and all functions , as well as merge_sort function

[maleadt]

debugging KernelAbstractions kernels is different from normal Julia

printf debugging? You can always do Main.OpenCL.@printf or something from KA.jl/AK.jl when you have OpenCL imported in your main environment (even though KA.jl/AK.jl doesn't depend on it, hence going through Main).

[vchuravy]

KA also has an @print

[VarLad]

I think the issue comes from me trying to use CPU as a GPU backend in AcceleratedKernels.jl
Also, I just saw JuliaGPU/KernelAbstractions.jl#556
I'll look into the issue again when the above PR lands :)

[vchuravy]

I think the issue comes from me trying to use CPU as a GPU backend in AcceleratedKernels.jl

Are you using CPU or OpenCLBackend

[VarLad]

OpenCLBackend with POCL on CPU

[vchuravy]

Okay, so that will be an issue for the default backend once JuliaGPU/KernelAbstractions.jl#556 lands.

[maleadt]

I think the issue comes from me trying to use CPU as a GPU backend in AcceleratedKernels.jl

I'm not sure what you mean by this. Every GPU construct in KA.jl should be supported by the OpenCL back-end, even when used on the CPU.

[VarLad]

I meant it as in, some algorithms are written with defaults according to GPU specifications in mind.
For example, I was able to make the functions in this file work by setting cooperation to false to use a less optimized implementation but one which gave correct results for POCL on CPU: https://github.com/JuliaGPU/AcceleratedKernels.jl/blob/main/src/truth.jl#L51
Without setting it to false, the functions any and all give incorrect outputs.

[maleadt]

That's unexpected; without looking into the implementation, non-UB code should execute equally well on an actual GPU as it does on PoCL's CPU back-end.

[VarLad]

Thats unexpected for me too, but the implementation is not working on specific GPUs as well (Intel UHD 620 integrated graphics cards and some other Intel GPUs as stated in the docstring I linked above), for which the flag exists. The flag is set to false for oneAPI because global writes hangs on some Intel GPUs apparently.

[VarLad]

Okay, I think I know the issue. Its probably in OpenCL.jl
Here's a reproducer:

using OpenCL, pocl_jll, KernelAbstractions

@kernel inbounds=true function _mwe!(@Const(v))
           temp = @localmem Int8 (1,)
           i = @index(Global, Linear)
           @print i "\n"
           @synchronize()
       end

v = CLArray(rand(Float32, 10))

_mwe!(OpenCLBackend(), 256)(v, ndrange=length(v))

This prints 1...256.

The CUDA version of the same code prints 1...10.

using CUDA

b = CuArray(rand(Float32, 10))

_mwe!(CUDABackend(false, false), 256)(b, ndrange=length(b))

The issue is probably that ndrange is not working.

Consequently, creating a CLArray of size 256 works without any issues, for the any and all functions , as well as merge_sort function.

[VarLad]

I think its multiple issues:

• merge_sort actually works on Floats of all sizes just fine, but Ints are an issue, where it replaces some elements by 0. In case of CLArray of Integers of size 256, I notice that the elements are not 0, but the sorting is still wrong. The specific issue still needs to be pinpointed, although I'm guessing that issue is something similar to the second point.
• ndrange not working how it should, is still an issue, irrespective of types, and is leading to wrong output (in case of any and all functions). Also, irrespective of types, sizes of 256 "just work" for the above two functions.

If the provided MWE doesn't make sense, or I misidentified the second issue, please let me know.

[maleadt]

Reduced to the common mistake of having barriers in divergent code:

function foobar(n)
    if get_local_id() <= n
        OpenCL.@show get_global_id()
        barrier()
    end
    return
end

@opencl local_size=16 global_size=16 foobar(10)

The problem is that the barrier is placed in a divergent section of code, which is undefined behavior. According to the OpenCL™ 1.2 specification, section 3.4.3:

[T]he work-group barrier must be encountered by all workitems of a work-group executing the kernel or by none at all.

Same applies to CUDA's _syncthreads/bar:

In conditionally executed code, a bar instruction should only be used if it is known that all threads evaluate the condition identically (the warp does not diverge).

@vchuravy This looks like a potential design flaw with KA.jl's implicit bounds checking?