OpenCL error

[peastman]

I'm porting an OpenCL application to run on Xeon Phi.  I've encountered what I'm pretty sure is a bug in the compiler or runtime.  With some work, I've narrowed it down to the following minimal test case.  To observe the problem, just launch testKernel() using a single workgroup of 64 work units.  It repeatedly prints out

1 <= 0????

void testFunction(__local short* a, __local int* b) {
   if (get_local_id(0) == 0)
       *a = 50;
   barrier(CLK_LOCAL_MEM_FENCE);
   int c = *a/32;
   if (c > 0) {
       if (get_local_id(0) == 0)
           *b = 0;
   }
   else
       printf("%d <= 0????\n", c);
}

__kernel void testKernel() {
   __local short a;
   __local int b;
   testFunction(&a, &b);
}

Any idea what the problem is, or how I can work around it?

Peter

[Yuri_K_Intel]

Hi, I will look into this issue. A couple of comments. It would be better (faster to reproduce) if you could provide full source code, including host part. There is a separate forum for OpenCL - http://software.intel.com/en-us/forums/intel-opencl-sdk, please use it for OpenCL related issues/questions. Thanks, Yuri

[peastman]

Hi Yuri,

Thanks for looking into it.  My contact at Intel specifically told me to post on this forum.  Do you want me to repost it on the other one?

I don't have a simple, self contained host program for launching it.  I'm working from a very large code base.  But all it really needs to do is the generic boilerplate found in any OpenCL program: select the platform, figure out which device is the Xeon Phi, create a context, compile the kernel, and launch it.

Peter

[Yuri_K_Intel]

Hi Peter, No need to repost. I was able to reproduce the issue and I have submitted a ticket for it. I will inform you of the further investigation results. Thanks, Yuri

[peastman]

Thank you!  Let me know what you find out, especially if you can suggest a workaround.

Peter