Open
Description
(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 CHIP-SPV/chipStar#215.
Environment:
- OS: Ubuntu 22.04.2 LTS
- DPC++ version: (not used directly)
- Intel OneAPI: 2023.1.0-46401 (intel-basekit APT)