Passing structs by value on an Intel GPU does not load struct members properly

[Zuzek__Peter]

Hi,

it looks like there is a bug in the Intel GPU OpenCL device: when passing
struct instances to a kernel by value (instead of a buffer object), enqueuing
the kernel (using clEnqueueNDRangeKernel) will raise error -52 (CL_INVALID_KERNEL_ARGS).
Of the tested devices, this only happens on the Intel GPU device and only with
some structs, not all. For example, the following:

struct s_int_short_int {
  int a;
  short b;
  int c;
};

get copied correctly, while the following:

struct s_short_char_char {
  short a;
  char b;
  char c;
};

fails while enqueuing.

I've provided a test case (attached to this post) that tries to copy various
structs either by value or by pointer (buffer object) on different devices
(Intel CPU and AMD GPU using their latest OpenCL drivers). Again, only the
Intel GPU device fails. The same behavior can be observed when using precompiled
SPIR instead of OpenCL C kernels.

This was tested on Windows 7 and on CentOS 7.2 using the Intel HD 530 GPU
- Windows reports OpenCL driver version 20.19.15.4463 (though we can also
reproduce it using 20.19.15.4474) and CentOS reports r2.0.54425.

Thank you,

Peter

-- 
Peter Žužek
Software Engineer
Codeplay Software Ltd
Level C Argyle House, 3 Lady Lawson Street, Edinburgh, EH3 9DR
Tel: 0131 466 0503
Fax: 0131 557 6600
Website: http://www.codeplay.com
Twitter: https://twitter.com/codeplaysoft

[Jeffrey_M_Intel1]

Hi Peter,

We've replicated here and agree that this definitely looks like a bug.  Thanks for bringing this finding to our attention with your excellent description and replicator.  I'll get this filed so we can start working on a fix as soon as possible.

Thank you,

Jeff