How to pass private array to OpenCL kernel as arguments?

Question

My kernel arguments are the following:

 __kernel void
codesGPU(struct stateGPU* s,
         short* lencnt,
         short* lensym,
         short* distcnt,
         short* distsym,
         __global const struct in_unit* input,
         __global struct out_unit* output,
         struct contextGPU* ctx,
         __global const short* lens,
         __global const short* lext,
         __global const short* dists,
         __global const short* dext,
         int* outsize
         )

input and output are "global" because they contain input/output data for all work-items, which are separated by get_global_id(0). lens, next, dists, dext are 4 constant global arrays shared by all work-items, so I declare them as global constant. However, the rest arguments are private to each work-item. These includes scalars like outsize and arrays like lencnt. I use pointers here because I still need to pass these arguments back to CPU after GPU computation.

My question is:

Is it possible to pass private arrays to kernels?
What's the best way to pass these variables (different for each work-item, and still want to return the value) to kernel?
How should I write the kernel code and C code?

The kernel also calls two helper functions in .cl code:

int fun1(struct stateGPU *s, short *countarr, short *symarr, __global const struct in_unit *input, int *idx, struct contextGPU* ctx, int did)

int fun2(struct stateGPU *s, int need, __global const struct in_unit* input, int *idx, struct contextGPU* ctx, int gid)

In the main.c, I create buffer objects for global variables and pass them as

cl_mem cl_lens;
cl_lens = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(lens), lens, &err);
err |= clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *)&cl_lens);

Can anyone give some suggestions on how to pass other arguments to the kernel, based on my problem? Many Thanks!

Don't you think these questions are just a bit general? Also, to whom is the array "private""? If the module to which the arrays are being passed does not enjoy access to those regions of memory, you will be in trouble. Always be acutely aware of segmentation register settings. — Bruce David Wilner

DarkZeros DarkZeros · Accepted Answer · 2016-07-14T08:22:10

Is it possible to pass private arrays to kernels?

No, only global/constant data can be passed to kernel. The only way to set a register is to pass it from somewhere, aka global memory.

What's the best way to pass these variables (different for each work-item, and still want to return the value) to kernel?

Simply global memory, since constant memory is more suited when the same data is needed for all work-items.

How should I write the kernel code and C code?

Well that depends on what you actually want it to do! :)

How to pass private array to OpenCL kernel as arguments?

1 Answers