send custom type to the kernel

Hi forum,

i need to send simulation parameters encapsulated inside a struct and i need to send the struct to the kernel.

What should be the kernel arguments type? Should it be as follows:


__kernel void ker(__global const float *a,
                  __global const float *b, 
                  __global float *c,
                  __constant struct Params* test
                  )

Thanks

If you have only one instance of the struct to send to the kernel, you can pass it by copy:

__kernel void ker(…, struct Params test)

I’ve wondered about this a few times. Does OpenCL really support passing derived data types (structs ) as kernel arguments? The C/C++ compiler and OpenCL compilers could choose to pad the storage of structure members differently (perhaps to support different architecture alignment requirements). So, I don’t think the straight approach will not work in general.

I have seen a suggestion to use the structure attribute((packed)) qualification in both the C/C++ host code declaration and the OpenCL code declaration. This would arrange all data members sequentially to one-byte boundaries increasing the likelihood the structure data layouts are compatible. Some description of packed here.

I didn’t realize until now that this type of attribute is included in the OpenCL standard as described here.

In a related question, I wonder if I should be care to specify an endianness attribute to every kernel argument supplied by the host?

One of the job of the OpenCL runtime is to marshal data between host and device transparently. Alignment and packing are defined in the OpenCL specification and are compatible with standard C usage (alignment on a power of 2 and so on…)

Alignment and packing attributes are present in OpenCL C so that you can exactly map host C structures to OpenCL C structures when the host structures use custom alignment and/or packing.

Similarly, if you pass an int to a kernel, OpenCL will handle endianness and transparently swap bytes if endianness is different on host and device.

Troubles happen if you cast a char* pointer to an int* pointer, or if you use reinterpreting cast such as:

int i;
char4 c = as_char4(i);

In this case, the result will be “implementation-defined”.

[QUOTE=utnapishtim;30255]One of the job of the OpenCL runtime is to marshal data between host and device transparently. Alignment and packing are defined in the OpenCL specification and are compatible with standard C usage (alignment on a power of 2 and so on…)

Alignment and packing attributes are present in OpenCL C so that you can exactly map host C structures to OpenCL C structures when the host structures use custom alignment and/or packing.

Similarly, if you pass an int to a kernel, OpenCL will handle endianness and transparently swap bytes if endianness is different on host and device.

Troubles happen if you cast a char* pointer to an int* pointer, or if you use reinterpreting cast such as:

int i;
char4 c = as_char4(i);

In this case, the result will be “implementation-defined”.[/QUOTE]

That is certainly encouraging.
Speculation about handling endianness is probably not an issue to have too much heartache over given the hardware implementations out there. But, I do wonder how this could be handled transparently?

Can I trust the compiler to understand and track the context for data types passed in and out of the kernel?


kernel void foo_kernel( global int * input_array )
{
     int native_opencl_int = input_array[0];  // might data type endianness conversion happen here automatically?

}

Or to ensure compatibility with the OpenCL kernel, do I need to use the exposed OpenCL API data types in my host code? e.g. cl_int?

Note that buffers use the endianness of the device, so a buffer should be read or written taking this into account.

You can change this behavior with attribute((endian(host))) to declare that a buffer is handled with host endianness.

What OpenCL handles automatically is the order of the components of a vector type: v.x is always the “first” component of a float4 vector, whatever the way it is stored in memory (unlike i[0][/i] for instance)

When arguments are passed by copy, OpenCL Specification states that “the implementation may or may not automatically convert endianness of kernel arguments. Developers should consult vendor documentation for guidance on how to handle kernel arguments in these situations.”