CODEWITHSUNDEEP
typesopenclgmp
Studer
Studer

Reputation: 641

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 ?

Upvotes: 14

Views: 19848

Answers (4)

gallo
gallo

Reputation: 594

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

Upvotes: 3

user2102611
user2102611

Reputation:

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 ♣

Upvotes: 8

VHristov
VHristov

Reputation: 1079

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.

Upvotes: 22

Adam S.
Adam S.

Reputation: 1271

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.

Upvotes: 4

Related Questions