Results 1 to 6 of 6

Thread: send custom type to the kernel

  1. #1

    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:

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

    Thanks

  2. #2
    Senior Member
    Join Date
    Oct 2012
    Posts
    105
    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)

  3. #3
    Junior Member
    Join Date
    Oct 2011
    Location
    Seattle, WA
    Posts
    27
    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?

  4. #4
    Senior Member
    Join Date
    Oct 2012
    Posts
    105
    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".

  5. #5
    Junior Member
    Join Date
    Oct 2011
    Location
    Seattle, WA
    Posts
    27
    Quote Originally Posted by utnapishtim View Post
    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".
    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?
    Code :
    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?

  6. #6
    Senior Member
    Join Date
    Oct 2012
    Posts
    105
    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 ((float*)&v)[0] 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."

Posting Permissions

  • You may not post new threads
  • You may not post replies
  • You may not post attachments
  • You may not edit your posts
  •