Complex Structs in openCL Kernel

Question

So I know that openCL has a lot of restrictions when it comes to what kind of arguments it can accept being passed into a kernel. For instance, it doesn't support C++ class types and a struct can't have a pointer as an argument. However, can openCL support structs that have non-pointer structs as fields?

For instance if I did the following:

typedef struct
{
    int a;
    float b;
} MyStruct1;

typedef struct
{
    float a;
    MyStruct1 field;
} MyStruct2;

__kernel void MyKernel(const MyStruct2 *struct)
{
    MyStruct1 innerStruct = struct->field;
    //Do more stuff
}

^^Would that be legal in openCL? And if it is legal, are there any special memory issues that I should be aware of?

You can try to copy two structs and debug if they do a deep copy or shallow one. If it is working for multiple of 128bit length structs, you should use those structs. — huseyin tugrul buyukisik, Oct 01 '14 at 20:35

score 3 · Answer 1 · edited May 23 '17 at 12:21

3

Nested structs are legal in OpenCL but padding is device and compiler specific. So it will lead to problems if use them as argument to kernel. On some systems padding will be different on host and on device which will lead to corrupted data. If you want to use structs on device side you can always move data with arrays and construct the structs on the device side.

edited May 23 '17 at 12:21

Community

1
1

answered Oct 01 '14 at 20:03

maZZZu

3,585
2
17
21

Ok, but is it legal for structs to contain other structs as fields? I know structs in general were legal but I'm wondering about this specific scenario of structs containing structs. – user1782677 Oct 01 '14 at 20:04
I meant nested structs when I was speaking of just structs. I already edited the answer. – maZZZu Oct 02 '14 at 05:23

score 2 · Accepted Answer · answered Oct 01 '14 at 20:06

2

OpenCL C, which is used inside kernels, is based on C99 standard with some restrictions. As soon as structure only describes position of it's elements in memory, there are no problems with nested structures.

OpenCL specs can be found at Khronos site

answered Oct 01 '14 at 20:06

Roman Arzumanyan

1,784
10
10

Complex Structs in openCL Kernel

2 Answers2