topic Comparison of get_global_id() to int differs when optimizations on/off in OpenCL* for CPU

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

Chris_C_1 — Mon, 31 Jul 2017 13:22:17 GMT

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

Thanks for this report. As I

Jeffrey_M_Intel1 — Mon, 31 Jul 2017 19:15:32 GMT

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?

Hi Jeffrey,

Chris_C_1 — Mon, 31 Jul 2017 21:39:00 GMT

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

The problem, as best as I can

Chris_C_1 — Mon, 31 Jul 2017 21:45:27 GMT

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.
  }
}

Thanks for this additional

Jeffrey_M_Intel1 — Wed, 02 Aug 2017 01:05:46 GMT

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?

Given the combination of

Chris_C_1 — Wed, 02 Aug 2017 09:55:27 GMT

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