Passing buffers (pointers) in a structure to a kernel

I have a kernel that use several ‘cl_mem’ buffers as input and outputs. They are currently passed using clSetKernelArgs. I will like to pass a single buffer that contains all the parameters for the kernel.

On the host side:

struct KernelArgs {
int mNumEntries;
cl_mem mTable;
cl_mem mResult;

on OCL side:
struct KernelArgs {
int mNumEntries;
global float* mTable;
global float* mResult;

__kernel void myKernel(
__const global KernelArgs * args


I cannot get this to work on OpenCL. Two problems:

  1. on AMD’s OCL I got the compile error: “kernel arguments can’t be declared with types”. No where in the 1.2 spec I can see this restriction.

  2. If I pass in ‘__constant char* argbuf’ instead and cast it to a typed ptr within the kernel. That worked for the scalar arguments - like ‘mNumEntries’. However, the pointers (which are cl_mem buffers) does not work. I verified I am getting the correct hex value for them by writing them back out into my output buffer. But using them to access the values in ‘mTable’ are returning garbage.

This works in CUDA and is a common practice. The alternative is to use a program scope constant buffer. Kind of like GLSL uniforms. But OpenCL program scope constant variables can only be initialized by compile time values.

Passing everything as kernel parameters is a little fragile since the host and gpu side code has to be carefully maintained.

If the above works I can use a shared header with a #define to pass arguments to the kernel.


This is not valid in OpenCL 1.x because the cl_mem pointer and the pointer the device code expects to see are not required to be the same (for reasons of portability). The runtime needs additional information to correctly map the pointer values into the device code and you cannot provide that information about the contents of your structure.