Avoiding data alignment in OpenCL

Question

I need to pass a complex data type to OpenCL as a buffer and I want (if possible) to avoid the buffer alignment.

In OpenCL I need to use two structures to differentiate the data passed in the buffer casting to them:

typedef struct
{
   char a;
   float2 position;
} s1;

typedef struct
{
   char a;
   float2 position;
   char b;
} s2;

I define the kernel in this way:

__kernel void 
Foo(
   __global const void* bufferData,
   const int amountElements // in the buffer
)
{
   // Now I cast to one of the structs depending on an extra value
   __global s1* x = (__global s1*)bufferData;

}

And it works well only when I align the data passed in the buffer.

The question is: Is there a way to use _attribute_ ((packed)) or _attribute_((aligned(1))) to avoid the alignment in data passed in the buffer?

Answer 1

If padding the smaller structure is not an option, I suggest passing another parameter to let your kernel function know what the type is - maybe just the size of the elements.

Since you have data types that are 9 and 10 bytes, it may be worth a try padding them both out to 12 bytes depending on how many of them you read within your kernel.

Something else you may be interested in is the extension: cl_khr_byte_addressable_store http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/cl_khr_byte_addressable_store.html

update: I didn't realize you were passing a mixed array, I thought It was uniform in type. If you want to track the type on a per-element basis, you should pass a list of the types (or codes). Using float2 on its own in bufferData would probably be faster as well.

__kernel void 
Foo(
   __global const float2* bufferData,
   __global const char* bufferTypes,
   const int amountElements // in the buffer
)