[maleadt]

Tracking PR for the CanonicalizeBarriers fix from pocl#2281, which upstream closed: they won't change the legacy work-item loop handlers and are moving kernel-compiler work to their loopvec rewrite instead. The JuliaGPU PoCL JLLs carry these commits as a local patch (Yggdrasil P/pocl: pocl_next 0004–0005, pocl 0004), so this PR keeps them rebased and reviewable until upstream main fixes these shapes another way.

---

Parallel-region formation treats the successors of a barrier block as entries of separate regions. That holds for conditional barriers, but not when a branch after the barrier reconverges before the next barrier: both successors then belong to one region, and Kernel::createParallelRegionBefore picks whichever comes first as its entry. IsolateRegions normally prevents this by giving the region a dedicated entry block, but only when LLVM's RegionInfo sees a single-entry single-exit region.

Two shapes from bounds-checked Julia kernels defeat that:

• One arm of the branch skips a variable-trip loop (the original reproducer, from a KernelAbstractions.jl radix sort). PoCL forms a malformed region ending at the loop latch and aborts in ParallelRegion::chainAfter.
• A guarded section between barriers has an early exit, e.g. the never-taken bounds-check failure path, that leaves the kernel without reaching the next barrier. Padding work-items branch from the barrier straight to the end of the section. The early exit makes the region multi-exit, so it gets no dedicated entry. Depending on the successor order, either the padding work-items enter the work-item loop past its per-work-item control and execute the guarded code (silently wrong results, also in release builds), or getParallelRegions forms a bogus region whose exit has two successors and WorkitemLoops::createLoopAround fails ExitBB->getTerminator()->getNumSuccessors() == 1. This is what KernelAbstractions.jl 0.10's trailing_uniform! and qualified! tests hit on PoCL main, through both the SPIR-V and the OpenCL C input paths (Implement KernelInterface, and support KernelAbstractions 0.10 OpenCL.jl#519).

The second shape only needs an early return/goto inside a guarded section, so it also occurs in plain OpenCL C:

kernel void skip_padding(global int *out, int n, int len) {
  int i = get_global_id(0);
  bool pad = i >= n;
  if (!pad) {
    if (i >= len) goto error;   // never taken
    out[i] += 1;
  }
  barrier(CLK_GLOBAL_MEM_FENCE);
  if (!pad) {
    if (i >= len) goto error;
    out[i] += 2;
  }
  barrier(CLK_GLOBAL_MEM_FENCE);
  return;
error:
  out[0] = -1;
}
// global size 16, local size 8, n = 13, len = 16:
// main with loops/loopvec: out[13..15] == 2 (expected 0)

The fix detects barrier successors that reach a common block without crossing another barrier, and moves their branch into a shared post-barrier block, as CanonicalizeBarriers already does for CBS. Branches whose arms end at distinct barriers keep the existing conditional-barrier handling.

Tests

• regression/barrier_divergent_skip: reduced SPIR-V for the loop-skipping shape.
• regression/barrier_early_exit: two OpenCL C kernels for the early-exit shape, one per symptom.

Both run with the loops, loopvec and CBS handlers. Without the fix, both fail with loops and loopvec (barrier_divergent_skip only with assertions enabled), and pass with CBS. The workgroup/, regression/ and kernel_compiler_unit_tests/ ctest subsets show no new failures (checked on macOS/arm64 with LLVM 20). KernelAbstractions.jl's test suite passes through OpenCL.jl on an asserts build with this change.