PoCL drops the stores of a work-group when a value lives across a barrier inside a branch #548

PoCL miscompiles a barrier inside a branch when a work-item value computed before the barrier is used after it: a work-group in which that value differs between work-items stores nothing, not even from the work-items for which it holds. The barrier is valid OpenCL, since all work-items of a work-group take the same branch.

using OpenCL, pocl_jll

# In the last work-group only, a per-work-item value is computed before a barrier and used
# after it. That work-group stores nothing.
src = """
kernel void k(global long *A, long n) {
  size_t I = get_global_id(0);
  if (get_group_id(0) == get_num_groups(0) - 1) {
    bool active = I < n;
    barrier(CLK_LOCAL_MEM_FENCE);
    if (active) A[I] = I + 1;
  } else {
    A[I] = I + 1;
  }
}
"""
prog = cl.Program(; source = src) |> cl.build!
kern = cl.Kernel(prog, "k")

n = 37
A = CLArray(zeros(Int, 40))
clcall(kern, Tuple{CLPtr{Int}, Int}, A, n; global_size = 40, local_size = 8)
@show Array(A)[1:n]

(Array(A))[1:n] = [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 0, 0, 0, 0, 0]

Work-items 32–36 of the last work-group should have stored 33–37. Variants of the kernel above:

Also wrong with PoCL main (b3033f8d3), so it isn't fixed upstream yet. OpenCL.jl 0.10.12, pocl_jll 7.2.2+1 (LLVM 22.1.8), Ryzen 9 9950X.

This came up in KernelAbstractions.jl, when compiling each @kernel body twice, without and with bounds checks, and choosing per work-group which copy runs (an alternative to JuliaGPU/KernelAbstractions.jl#849). With @synchronize that puts a barrier under a group-dependent branch, and the partial last work-group then stored nothing.