Custom types in OpenCL kernel

Question

Is it possible to use custom types in OpenCL kernel like gmp types (mpz_t, mpq_t, …) ?

To have something like this (this kernel doesn't build just because of #include <gmp.h>) :

#include <gmp.h>
__kernel square(
   __global mpz_t* input,
   __global mpz_t number,
   __global int* output,
   const unsigned int count)
{
   int i = get_global_id(0);
   if(i < count)
       output[i] = mpz_divisible_p(number,input[i]);
}

Maybe by adding different arguments to the fourth parameter (options) of clBuildProgram ?

Or does OpenCL already have types that can handle large numbers ?

Answer 1

You can use custom types but anything used in the kernel needs to be specifically written for OpenCL. Check out this website perhaps for how to implement larger precision numbers: FP128

Edit: NVIDIA's CUDA SDK has a complex number data type, it's not ideal but may give you some ideas on how they go about it, OpenCL should be similar.

Answer 2

Generally you can use any types in an OpenCL program. But since imports do not work, you have to re-define them within the same program. For example:

typedef char my_char[8];

typedef struct tag_my_struct
{
    long int        id;
    my_char         chars[2];
    int             numerics[4]
    float           decimals[4];
} my_struct;

__kernel void foo(__global my_struct * input,
                  __global int * output)
{
    int gid = get_global_id(0);
    output[gid] = input[gid].numerics[3]== 2 ? 1 : 0;
}

However, you obviously need to keep the definitions within and outside OpenCL the same. Also make sure the type has the same size on both device and host (using a sizeof(my_struct) should do the trick). In some cases I had to adjust the definitions, to have matching sizes.