can we use structure in opencl

hi forum,
can we use structure in opencl as we do in C language.if yes how can i do this

thanks

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.

can u give me the example pls…

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…)

You can, but you have to be careful that structure alignment matches between host and GPU

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.

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).

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

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.

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.