Passing a kernel global buffer to a local function in OpenCL

Question

Say we have a kernel function:

kernel void function(global const float* a, global const float* b, global float* c, int nElements)
{
...
c[gid] = a[gid] * b[gid];
}

but want to break up a large complex kernel into several smaller functions. How do I pass the global buffers to these smaller functions?

If I do the following I get an error of the form "implicit declaration of function 'cl_axpbyr' is invalid in OpenCL":

kernel void function(global const float* a, global const float* b, global float* c, int     nElements)
{
...
cl_axpbyr(1.0f, a, c, nElements);
}

inline void cl_axpy(float alpha, global const float* x, global float* y, int nElements)
{
int gid = get_global_id(0);
if (gid >= nElements)
 {
      return;
 } 
y[gid] = alpha*x[gid] + y[gid];
}

If the answer satisfies all you needs you should close the question to keep it from the list of unanswered questions. Otherwise it will clutter the list and take space from questions that actually require answering. — sharpneli

sharpneli sharpneli · Accepted Answer · 2013-11-11T11:15:08

First of all you call this:

cl_axpbyr(1.0f, a, c, nElements);

While your function is:

inline void cl_axpy

You should call cl_axpy instead of cl_axpbyr

Second of all OpenCL kernel language is just C. So you need to predeclare your functions if you are going to define them after the place you are going to call them. The following code compiles cleanly:

// This is the normal C style function declaration which must exist
inline void cl_axpy(float alpha, global const float* x, global float* y, int nElements);


kernel void function(global const float* a, global const float* b, global float* c, int     nElements)
{
cl_axpy(1.0f, a, c, nElements);
}


inline void cl_axpy(float alpha, global const float* x, global float* y, int nElements)
{
int gid = get_global_id(0);
if (gid >= nElements)
  {
      return;
  } 
y[gid] = alpha*x[gid] + y[gid];
}

You could also just place the whole cl_axpy before your kernel definition. Both ways work fine.

Passing a kernel global buffer to a local function in OpenCL

1 Answers