OpenCL rotate produces wrong results

dreifachstein commented 3 years ago

System: Debian unstable intel-opencl-icd 21.32.20609-1

Processor: i5-11300H

Kernel:

__kernel
void test_rotate32(
    __global uint* dst
)
{
  uint i = get_global_id(0);
  dst[i] = rotate(1U, i);
}

Executing the kernel for 0-31 on produces:

0000000000000001
0000000000000002
0000000000000004
0000000000000008
0000000000000010
0000000000000020
0000000000000040
0000000000000080
0000000000000100
0000000000000200
0000000000000400
0000000000000800
0000000000001000
0000000000002000
0000000000004000
00000000ffff8000
0000000000000001
0000000000000002
0000000000000004
0000000000000008
0000000000000010
0000000000000020
0000000000000040
0000000000000080
0000000000000100
0000000000000200
0000000000000400
0000000000000800
0000000000001000
0000000000002000
0000000000004000
00000000ffff8000

HoppeMateusz commented 3 years ago

this seems to be an issue in the compiler https://github.com/intel/intel-graphics-compiler , Looks like rotate operates on short type - and overflows after 00004000 to ffff8000.

JacekDanecki commented 3 years ago

Transferring to IGC project

mnaczk commented 2 years ago

Issue resolved by commit https://github.com/intel/intel-graphics-compiler/commit/78db6ad5b66839715665bd44bef745234d354511 Thanks @dreifachstein for detecting the problem. Thanks to @Quetzonarch and @jgu222 for the solution.

intel / intel-graphics-compiler

OpenCL rotate produces wrong results #211