View Full Version : can we use structure in opencl

12-05-2012, 03:56 AM
hi forum,
can we use structure in opencl as we do in C language.if yes how can i do this


12-05-2012, 02:01 PM
I use one big one; it holds all of my input parameters.

I pass in its pointer as my kernel arg 0, and then just dereference as usual to get to all of the elements. You can either £include the header file that defines it (with absolute path), or copy the definition into the top of your .cl file.

12-05-2012, 06:46 PM
can u give me the example pls......

12-06-2012, 05:30 PM
Well, good sir, I won't post my actual code for trade-secret reasons, but I can describe an example....

There is a structure "varholder" that holds everything the kernel needs to know. Simplified version:

struct hoohah
float xc[4][3], yc[4][3];
float zoom, twirl, funky, monkey;
int sizex, sizey, sizdv;
char Invert[4][3];
char thing0, thing1, thing2, thing3;
int kyoffset;
char frizzle[401];
typedef struct hoohah varholder;

... right now I have this inside my .cl file, but it also works to #include a header file, as long as I specify the full pathname:

#include "/Developer/velprojs/thisproj/kernelischestuff.h"

... then, the declaration of the kernel function itself looks like this:

__kernel void DoThisStuff( __global varholder * vh, __global __write_only uint * obuf )

... and then, everything in the structure is available by dereferencing through that pointer, i.e. vh->sizex, vh->thing0, et cetera....

-- BUT --

what I think you *cannot* do is have pointers in your structure. (at least you can't pass pointers into your kernel; maybe once you're inside you can make a structure that contains pointers, but i don't know. maybe somebody else can answer that....)

12-08-2012, 11:34 AM
You can, but you have to be careful that structure alignment matches between host and GPU

12-08-2012, 08:04 PM
Thanks, Dm --

Yes, there is one place in my structure, after lots of chars, that I inserted "char dum00;" so that all the following floats and ints would start on 4-byte boundaries. I had forgotten about that.

12-08-2012, 09:36 PM
thanks for replies i am asking that
1) if i make structure in c and pass there base address
in clcreatebuffer as i do in case of array ...is it possible or any problem

12-09-2012, 06:59 AM
thanks for replies i am asking that
1) if i make structure in c and pass there base address
in clcreatebuffer as i do in case of array ...is it possible or any problem

Right, that's how you would do it (just like an array).

12-12-2012, 11:09 PM
cpu and gpu i think, or no?
Or is possible mantain the same work items and works groups
Is possible to use other tools of intel like a prof

12-23-2012, 05:53 AM
Just a note for getting the alignment right, I think you should be using the cl_* types, i.e. cl_char, cl_float, etc. for the fields in your structure. Looking at cl_platform.h shows me that alignment attributes are specified for some of those types.

12-24-2012, 06:13 PM
One important thing to keep in mind is also that size_t and ptrdiff_t are likely not to be the same size on the host and on the GPU, which can cause misalignment

//Host code
struct my_struct{
int* my_ptr;
int a;

//Kernel code
struct my_kernel_struct{
int* my_ptr;
int a;

Here, a raw copy between a my_struct structure and a my_kernel_stuct structure might not give the expected result if your host is on a 64bit OS and your device works with 32bit arithmetic, which is often the case.
If portability is not an issue, you can querry CL_DEVICE_ADDRESS_BITS , and check that this is the same as your host.