Array in struct bug on Fermi

I recently got a GTX 480, which replaced my GTX 285, and updated to version 3.1 of the SDK. I have since then had problems with structures containing arrays. I have attached a small example that demonstrates the problem, and the following lines of CL kernel code is an excerpt of that. If run, the next cl-call returns CL_INVALID_COMMAND_QUEUE.

typedef struct Data


  float data[1];

} Data;

float get(Data data) {



__kernel void bug_kernel(__global Data* datas)


  Data data0 = datas[0];

  Data data1 = datas[1];[0] = get(data1);

  datas[0] = data0;


Have anyone else had similar problems? Does the attached application give other results on your machines? (4.45 KB)

Technically, composite types like structs and pointers which are nested cannot be passed as kernel arguments, so the behavior is undefined…

I am aware that pointers in structs are problematic, but pointers and arrays are not the same thing and I cannot remember anyone pointing out that arrays in structs should produce undefined behavior. Is this explicitly stated in the specification, or any other documentation?