Compiler error with Intel HD

[jiri]

Hi,

I'm having trouble compiling a kernel for Intel HD Graphics (either using OpenCL API or Code Builder). Compilation works for CPU. The kernel looks like this:

#define ACC_SIZE 32
__kernel void go(__global int* data, int iterations)
{ 
  int acc[ACC_SIZE];
  for(int i=0;i<ACC_SIZE;++i) acc=data[get_global_id(0)+i];//not necessary to break the compilation
  for(int i=1;i<iterations;++i) {
    for(int i=0;i<ACC_SIZE;++i) acc*=i;
  } 
  int sum=0;
  for(int i=0;i<ACC_SIZE;++i) sum+=acc;//not necessary to break the compilation
  data[get_global_id(0)]=sum;
}

It works for ACC_SIZE up to 16, but does not work for larger values. The code is not supposed to do some useful work, it is just a little experiment of mine. I get the following error message:

fcl build 1 succeeded.
error: Cannot select: 0x30ce758: i64,glue = adde 0x2e2e270, 0x2e2e270, 0x30ce6d0:1 [ID=117]
  0x2e2e270: i64,ch = CopyFromReg 0x2e2d638:1, 0x2e2dc10 [ORD=311] [ID=111]
    0x2e2dc10: i64 = Register %vreg49 [ORD=311] [ID=2]
  0x2e2e270: i64,ch = CopyFromReg 0x2e2d638:1, 0x2e2dc10 [ORD=311] [ID=111]
    0x2e2dc10: i64 = Register %vreg49 [ORD=311] [ID=2]
  0x30ce6d0: i64,glue = addc 0x2e2d638, 0x2e2d638 [ID=110]
    0x2e2d638: i64,ch = CopyFromReg 0x2dd19f8, 0x2e2ceb0 [ORD=311] [ID=105]
      0x2e2ceb0: i64 = Register %vreg48 [ORD=311] [ID=1]
    0x2e2d638: i64,ch = CopyFromReg 0x2dd19f8, 0x2e2ceb0 [ORD=311] [ID=105]
      0x2e2ceb0: i64 = Register %vreg48 [ORD=311] [ID=1]
In function: go
error: midlevel compiler failed build. 

Could someone either check whether it works for them or point me to some error in the code?

Thanks.

[Robert_I_Intel]

Hi Jiri,

Please upgrade your system to use the latest Intel(R) Processor Graphics drivers. We also just released a new version of Intel(R) INDE, which has an updated version of OpenCL(TM) Code Builder. Please follow the steps outlined in my article here: https://software.intel.com/en-us/articles/getting-started-with-opencl-development-on-windows-with-intel-inde if you just want to install OpenCL(TM) related tools.

I am unable to reproduce your error on my machine, which has the latest drivers and tools installed. If after updating the drivers and tools you still have the issue, please let me know what processor you have, what OS you are running on, and what versions of drivers and OpenCL tools you use.

[jiri]

Hi,

thank you for your quick reply. I already installed the latest INDE before submitting the original message (Code Builder shows version 1.3.0.43). Since my HD4000 is in a Lenovo notebook (the CPU is i5-3210M), I have trouble installing the latest drivers, since Lenovo also provide their own and is seems the drivers are not that happy with each other. However, based on your message, which suggest it is probably driver related (as I think the compiler is in the drivers, not in the Code Builder), I decided to try a bit more and do the update. It took some playing around, but now I managed to update to Intel-provided drivers. The drivers on the "OpenCL™ Drivers and Runtimes for Intel® Architecture" page do not seem to support HD4000, but I found a different driver ( win64_153332.zip ), which is also new ( 01/08/2015 ), but supports HD4000. After update, windows (I run Win7 64bit) say that it is version 10.18.10.4061 from 18.12.2014, which sounds reasonable to me. The clang_compiler64.dll, which could be relevant to the issue, is said to be version 3.0.1.10878.

[Robert_I_Intel]

Hi Jiri,

I was able to replicate the issue on my machine with HD4000: thanks for reporting the issue - I will file a bug with our OpenCL driver team.

[jiri]

Thanks. The code was a work-in-progress for an experiment I want to make and the subsequent versions are no longer affected by this issue, so the fix is not critical for me. But since it could surface somewhere else, possibly with someone's real application, I wanted to check whether that was a real bug or just problem with my system or code.

This is not the first time I managed to "break" a compiler used by Intel, but before it was always Intel C++ compiler when I played with offloading (very early implementations) to the Xeon Phi. I also once "scored a hit" on GCC, by mixing templates and the ?: ternary operator, so my "skills" are not limited to Intel software. It seems I'm not that good at using tools the way the authors anticipated (and tested) :-). Anyway, it's good to see that such positive response and desire to fix things is the norm at Intel.