CODEWITHSUNDEEP
opencl
Graham Seed
Graham Seed

Reputation: 852

Passing a kernel global buffer to a local function in OpenCL

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];
}

Upvotes: 0

Views: 1155

Answers (1)

sharpneli
sharpneli

Reputation: 1621

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.

Upvotes: 1

Related Questions