[OpenCL CPU Runtime] Kernel sensitive failure?

[linehill]

(Not sure if this is the correct place for this issue)

I have two SPIR-V module, a working one nobug.spv which has following kernels in pseudo OpenCL:

global int *Foo;
kernel void initFoo(global int *Ptr) { Foo = Ptr; }
kernel void readFoo(global int *Out) { *Out = *Foo; }

And failing one bug.spv which is the same but readFoo is defined as (in pseudo code):

//…
kernel void readFoo(global int *Out) { 
int tid = get_local_id(0) + get_group_idx(0) * get_local_size(0);
if (tid == 0)
  Out[tid] = *Foo; 
}

The attached OpenCL test application bug.cpp fails with bug.spv but passes with nobug.spv on Intel OpenCL CPU:

$ source /opt/intel/oneapi/setvars.sh
<…>
$ c++ -Wall bug.cpp -lOpenCL -o repro
$ clinfo -l
Platform #0: Intel(R) FPGA Emulation Platform for OpenCL(TM)
 `-- Device #0: Intel(R) FPGA Emulation Device
Platform #1: Intel(R) OpenCL
 `-- Device #0: Intel(R) Core(TM) i5-6400 CPU @ 2.70GHz
Platform #2: Intel(R) OpenCL HD Graphics
 `-- Device #0: Intel(R) HD Graphics 530
Platform #3: Portable Computing Language
 `-- Device #0: pthread-Intel(R) Core(TM) i5-6400 CPU @ 2.70GHz
$  ./repro 1 0 nobug.spv
Platform: Intel(R) OpenCL
  Device: Intel(R) Core(TM) i5-6400 CPU @ 2.70GHz
PASSED!
$ ./repro 2 0 bug.spv
Platform: Intel(R) OpenCL HD Graphics
  Device: Intel(R) HD Graphics 530
PASSED!
$ ./repro 3 0 bug.spv
Platform: Portable Computing Language
  Device: pthread-Intel(R) Core(TM) i5-6400 CPU @ 2.70GHz
PASSED!
$ ./repro 1 0 bug.spv
Platform: Intel(R) OpenCL
  Device: Intel(R) Core(TM) i5-6400 CPU @ 2.70GHz
Error: ReadVal: expected '-456'. Got '-478245456'

The issue is affected by https://github.com/CHIP-SPV/chip-spv/issues/215.

Environment:

• OS: Ubuntu 22.04.2 LTS
• DPC++ version: (not used directly)
• Intel OneAPI: 2023.1.0-46401 (intel-basekit APT)

repro.zip

[AllanZyne]

Hi @linehill, I think the if-condition should be tid == 0 in bug.spv, right?

[linehill]

I think the if-condition should be tid == 0 in bug.spv, right?

Yes, that's correct. I fixed the description.

[AllanZyne]

A patch for this bug has been submitted into the internal repo, please wait for next release. Thanks for catching this subtle bug!