Comparison of get_global_id() to int differs when optimizations on/off

[Chris_C_1]

Hi,

With the following kernel:

__kernel void A(__global double* input) {
  int tid = get_global_id(0);
  if (tid < - 1) {
    input[tid] = 1;
  }
}

The 'if' branch of the conditional should never be executed, since global id is always >= 0. However, on my E5-2620 and i5-4570, it is. When optimizations are disabled, it isn't. Could this be a compiler bug?

Full test case is here: http://paste.ubuntu.com/25213218/ 

My driver version is 1.2.0.25.

Cheers,
Chris

[Jeffrey_M_Intel1]

Thanks for this report.  As I've investigated, I have not been able to reproduce tid having a value of <0 with or without optimizations, with the CPU-only opencl-1.2-6.4.0.25 release or with SRB5.  Could you tell us a bit more about your system configuration, and the results you are seeing?   The contents of /etc/OpenCL/vendors could help us make sure we are testing the same OpenCL version.  What do you see there?

[Chris_C_1]

Hi Jeffrey,

Thanks for the quick reply. My driver version is:

$ cat /etc/OpenCL/vendors/intel64.icd
/opt/intel/opencl-1.2-6.4.0.25/lib64/libintelocl.so

The OpenCL devices on which I've observed the incorrect output are "Intel(R) Xeon(R) CPU E5-2620 v4 @ 2.10GHz" and "Intel(R) Core(TM) i5-4570 CPU @ 3.20GHz", and the host OS is Ubuntu 16.04.

With the test case I linked above, the output is consistently:

$ gcc foo.c -lOpenCL && ./a.out 2>/dev/null
computed result: 1.000
$ gcc foo.c -DNO_OPT -lOpenCL && ./a.out 2>/dev/null
computed result: 0.000

On all other devices I've checked (including an E5-2650 running opencl-1.2-4.4.0.117 release), the output is consistently 0.000.

Cheers,
Chris

[Chris_C_1]

The problem, as best as I can tell, appears to be in the implicit conversion of get_global_id() from size_t to int. On my system I am unable to coax the comparison of get_global_id() into int type. This kernel has the equivalent (incorrect) behavior when optimizations are enabled:

__kernel void A(__global double* input) {
  if (((int)get_global_id(0)) < ((int)- 1)) {
    input[tid] = 1; // Executed when opt enabled, not executed when -cl-opt-disable.
  }
}

[Jeffrey_M_Intel1]

Thanks for this additional info.  It is interesting that this is only happening on a few types of machines.  I'm still getting a 4th Generation Core machine set up for more investigation with the goal of filing a bug.  

In the meantime, the team is wondering about the context of how you found this behavior -- in validation of a larger group of kernels, or in new development?

Also, you mentioned that  you have an e5-2650 that does not have the problem.  Is this also v4?  If you use a different OpenCL runtime release on your E5-2620 v4 or  i5-4570 do you see the expected results?

[Chris_C_1]

Given the combination of drivers and hardware I have, I can't be sure which is at fault. I've been able to replicate this bug on:

• "Intel(R) Xeon(R) CPU E5-2620 v4 @ 2.10GHz" with driver "1.2.0.25" running Ubuntu 16.04
• "Intel(R) Core(TM) i5-4570 CPU @ 3.20GHz" with driver "1.2.0.25" running Ubuntu 16.04

But not on:

• "Intel(R) Xeon(R) CPU E5-2650 v2 @ 2.60GHz" with driver "1.2.0.44" running CentOS 7.1

The discrepancy was found during batch testing a group of kernels, though I made sure to manually reproduce on the E5-2620.

Hope that helps.

Cheers,
Chris