Array in struct bug on Fermi

Accelerated Computing CUDA CUDA Programming and Performance
1

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) {

  return data.data[0];

}

__kernel void bug_kernel(__global Data* datas)

{

  Data data0 = datas[0];

  Data data1 = datas[1];

  data0.data[0] = get(data1);

  datas[0] = data0;

}

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

2

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

3

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?