Custom types in OpenCL kernel

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 ?

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.

I used VHristov's answer and dietr's comment to get mine working. This code works for me in OpenCL 1.2

kernel

typedef struct tag_my_struct{
  int a;
  char b;
}my_struct;

__kernel void myKernel(__global my_struct *myStruct)
{
    int gid = get_global_id(0);
    (myStruct+gid)->a = gid;
    (myStruct+gid)->b = gid + 1;
}

host

typedef struct tag_my_struct{
  cl_int a;
  cl_char b;
}my_struct;

void runCode() 
{
    cl_int status = 0;
    my_struct* ms = new my_struct[5];

    cl_mem mem = clCreateBuffer(*context, 0, sizeof(my_struct)*5, NULL, &status);
    clEnqueueWriteBuffer(*queue, mem, CL_TRUE, 0, sizeof(my_struct)*5, &ms, 0, NULL, NULL);

    status = clSetKernelArg(*kernel, 0, sizeof(ms), &mem);

    size_t global[] = {5};
    status = clEnqueueNDRangeKernel(*queue, *kernel, 1, NULL, global, NULL, 0, NULL, NULL);

    status = clEnqueueReadBuffer(*queue, mem, CL_TRUE, 0, sizeof(my_struct)*5, ms, 0, NULL, NULL);

    for(int i = 0; i < 5; i++)
        cout << (ms+i)->a << " " << (ms+i)->b << endl;
}

output

0 ☺

1 ☻

2 ♥

3 ♦

4 ♣

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.

If you want to include header files into kernel file, you can add the -l dir as an argument to clBuildProgram, where dir is the directory with the header files.

More explanation here: include headers to OpenCL .cl file

Source: http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clBuildProgram.html