Hacktoberfest 2026: the issues maintainers tagged for October, open and beginner-friendly. Browse Hacktoberfest issues

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

Open
  • Is something wrong?
#548 0 comments 0 reactions 0 assignees View on GitHub

Maintainers usually reply within 1 day

Nobody has claimed this yet.

Assessment

Difficulty
4/5
Estimated time
3-5 days
Newbie friendliness
43/100
Issue type
Bug
Clarity
Clearly specified
Activity status
Active
Tech stack
julia
Domain
backend, testing-qa

Research direction

Start with the reproducer in the issue and confirm the failure using the stated OpenCL.jl and PoCL versions; compare the listed variants to isolate the branch/barrier behavior. The payload names no source files or tests, so locate the PoCL compiler path responsible for this behavior and its existing regression tests. Done when the reproducer stores the expected values and a regression test covers it.

Written by the indexing model from the issue text.

Description

bug upstream

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:

variant result
as above wrong
if (n > 0) around the barrier (all work-groups take it): the last work-group still stores nothing wrong
int instead of long/size_t wrong
I < n evaluated after the barrier instead of before it correct
no branch around the barrier correct
POCL_WORK_GROUP_METHOD=cbs correct
built with -cl-opt-disable correct

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.

Dominant language
Julia
Stars
313
Forks
50
Avg merge
23h 33m
Merged PRs (30d)
35

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

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

More from JuliaGPU/OpenCL.jl

All issues in JuliaGPU/OpenCL.jl

Similar issues

More Julia issues

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.