Wrong results from a 10-line OpenCL kernel on the Intel OpenCL CPU runtime (optimized build only)

[MarkRose]

Describe the bug

A kernel that updates a pair of private doubles with a butterfly step, a' = a + c*b, b' = a - c*b, once before and once after a work-group barrier, produces wrong values for b on the Intel OpenCL CPU runtime when the program is built at the default optimization level. The same source built with -cl-opt-disable on the same device is correct, and Intel's GPU compiler (NEO, Intel UHD Graphics) and POCL 7.2 agree with that result bit for bit. The failure does not depend on the vectorizer: with CL_CONFIG_USE_VECTORIZER=False the result is still wrong, just differently.

Small changes make the optimized build correct, which points at the optimizer rather than the arithmetic:

• removing the barrier, or making the second update a plain a = a + b;
• writing the product once, p = c * b; a = t + p; b = t - p; instead of a = fma(c, b, t); b = fma(-c, b, t); (the plain-arithmetic form a = c*b + t; b = t - c*b; still fails);
• dropping the negation (b = fma(c, b, t)) also makes it pass.

Only b is wrong, and only for work-items 0..55 of the 64; work-items 56..63 are correct (see the analysis under Additional context).

This was found in PRPLL (https://github.com/gwoltman/gpuowl), a Mersenne-prime tester, where the FFT kernels contain exactly this pattern and the program's Gerbicz error check fails on the CPU runtime while every GPU passes.

To reproduce

1. Kernel (repro.cl):

kernel __attribute__((reqd_work_group_size(64, 1, 1)))
void repro(global double2 * restrict out, const global double2 * restrict in, global const double2 * restrict trig) {
  uint me = get_local_id(0);
  double a = in[me].x, b = in[64 + me].x;
  double c = trig[me].x;
  { double t = a; a = fma(c, b, t); b = fma(-c, b, t); }
  barrier(CLK_LOCAL_MEM_FENCE);
  { double t = a; a = fma(c, b, t); b = fma(-c, b, t); }
  out[me].x = a;
  out[64 + me].x = b;
}

repro.c

The attached host program repro.c builds this kernel twice with clBuildProgram, once with the options given on the command line and once with -cl-opt-disable appended, runs each once on a single work-group of 64 work-items with deterministic pseudo-random inputs in [-1, 1), and compares the two outputs. REPRO_DUMP=1 prints every value.

2. Compile:

gcc -O1 -o repro repro.c -lOpenCL -lm

3. Run (<device-index> counts devices across all platforms in clGetPlatformIDs order; ./repro 0 repro.cl prints the device name first):

./repro <device-index> repro.cl -cl-std=CL2.0

Also fails with -cl-std=CL3.0, and with CL_CONFIG_USE_VECTORIZER=False ./repro ....

4. What is wrong and what was expected:

Intel OpenCL CPU runtime, Intel(R) Core(TM) i9-10885H CPU @ 2.40GHz (OpenCL 3.0 (Build 0)):

optimized      : out[0] = (1.4467647533780621, 0)  checksum -21.61488150677485
-cl-opt-disable: out[0] = (1.4467647533780621, 0)  checksum -16.01809545024355
optimized vs -cl-opt-disable: 56 of 256 doubles differ (max rel diff 87.9)
first difference: double2 index 64 (.x): optimized 1.4691584900322709 vs opt-disable 1.5436134085142328
differing double2 indices: 64 65 66 ... 119
RESULT: MISMATCH

Expected: the two builds agree (exit code 0), as they do on the other implementations tried on the same machine:

device: Intel(R) UHD Graphics (OpenCL 3.0 NEO )                     optimized = -cl-opt-disable, checksum -16.01809545024355, RESULT: match
device: cpu-haswell-... (OpenCL 3.0 PoCL 7.2, LLVM 21)               optimized = -cl-opt-disable, checksum -16.01809545024355, RESULT: match

The -cl-opt-disable result on the CPU runtime has the same checksum as both of those, so it is the correct one. A straightforward evaluation of the two butterfly steps in Python (with math.fma) also reproduces the -cl-opt-disable values exactly.

Environment

• OS: Linux. Ubuntu 24.04 (x86_64), running as a container on a Linux 6.18 host. The failure was first seen on the same host.
• Target device and vendor: Intel CPU, Intel(R) Core(TM) i9-10885H (Comet Lake, AVX2, no AVX-512), through the Intel OpenCL CPU runtime. Platform Intel(R) OpenCL, device version OpenCL 3.0 (Build 0), driver version 2026.21.7.0.24_160000.
• DPC++ version: no DPC++ compiler installed, runtime packages only, from https://apt.repos.intel.com/oneapi: intel-oneapi-runtime-opencl 2026.1.1-325, intel-oneapi-runtime-compilers 2026.1.1-325, intel-oneapi-runtime-dpcpp-sycl-opencl-cpu 2026.1.1-325, intel-oneapi-runtime-tbb 2023.1.0-151. ICD loader: ocl-icd-libopencl1 2.3.2 (Ubuntu). The OpenCL program is plain OpenCL C compiled by the runtime; no SYCL involved.
• Dependencies version: sycl-ls is not part of the runtime packages; the full clinfo output is attached (clinfo-intel-cpu.txt).

clinfo-intel-cpu.txt

Additional context

Where the wrong values come from: comparing the optimized output with an exact evaluation of the two butterfly steps shows that a is correct for every work-item, and that in the second step b of work-item m was computed with the multiplier c of work-item 56 + (m mod 8) instead of its own, for every m < 56:

m -> work-item whose c was used for b after the barrier
0:56 1:57 2:58 3:59 4:60 5:61 6:62 7:63 8:56 9:57 ... 55:63 56:56 57:57 ... 63:63

So the value that reaches the second fma(-c, b, t) after the barrier is the one belonging to the last group of 8 work-items (56..63) for everybody, which is why those 8 are correct. c itself is fine, only the negated copy -c is affected. The optimized code seems to compute -c once before the barrier for both uses, and then not save and restore that value per work-item across the barrier, so after the barrier every work-item reads the last vector of 8 that was computed. This explains all the variations above: with a single fma(c, b, t) after the barrier (no -c in the second step) or with the product written once (p = c * b) there is no such shared value, and the result is right. With the vectorizer off, the same mechanism gives a different wrong result (work-item 63's value, presumably).

The reproducer was reduced mechanically from PRPLL's tailMul kernel (src/cl/tailmul.cl and src/cl/fftbase.cl, height-FFT variant 2); in the original code the pair update is the X2_via_FMA macro and the value c is one of six trig values preloaded into a private array before the barrier. In the full program, -cl-opt-disable, marking any of the three helper functions noinline, or declaring the array volatile all give correct results; rewriting the fma() calls does not.

[MarkRose]

@againull Further reduction: it doesn't need fma or the butterfly at all. It happens whenever the result of an fneg is live across a barrier:

kernel __attribute__((reqd_work_group_size(64, 1, 1)))
void repro(global double2 * restrict out, const global double2 * restrict in, global const double2 * restrict trig) {
  uint me = get_local_id(0);
  double n = -trig[me].x;
  barrier(CLK_LOCAL_MEM_FENCE);
  out[me].x = n;
  out[64 + me].x = 0;
}

This gives 63 of 64 wrong with CL_CONFIG_USE_VECTORIZER=False and 56 wrong with the vectorizer on. The work-item-63 guess in the report is correct: with the vectorizer off, the asm has no per-work-item save of the value. After the first work-item loop it loads trig[63], negates it once and stores that for every work-item.

It depends on the IR opcode, not on negation as such:

• Fails: -c on double, float or double2, -(double)get_local_id(0), and -1.0 * c (the optimizer turns it into fneg). It also fails with CLK_GLOBAL_MEM_FENCE.
• Passes: 0.0 - c (fsub), 3.0 * c, fabs(c), integer negation, copysign(c, -c), and as_double(as_ulong(c) ^ 0x8000000000000000UL). It also passes when the value being negated is the same for every work-item.

So the code that decides which values differ per work-item probably doesn't treat fneg (a unary operator, not a binary one) as differing, and the value never gets saved per work-item across the barrier. Tested with runtime 2026.1.1, driver 2026.21.7.0.24_160000

[againull]

@MarkRose Yes, thank you, it's indeed the root cause. I merged the fix into corresponding repository. I will let you know the version that contains the fix and when it is updated here: https://github.com/intel/llvm/blame/sycl/devops/dependencies.json#L36