Call Kernelfunction from Other Kernel

Hi,

I found the Lines “It is just a regular function call if a __kernel function is called by another kernel function”, but if I try something like this,

[codebox]__kernel float

five()

{

return 5.0;

}

__kernel void

add five(__global float *a, __global float answer)

{

int gid = get_global_id(0);

answer[gid] = a[gid] + five();

}[/codebox]

this

[codebox]__kernel __local float

five()

{

return 5.0;

}

__kernel void

add five(__global float *a, __global float answer)

{

int gid = get_global_id(0);

answer[gid] = a[gid] + five();

}[/codebox]

this

[codebox]__kernel

five(__local float &val)

{

val = 5.0;

}

__kernel void

add five(__global float *a, __global float answer)

{

int gid = get_global_id(0);

__local val = 0;

five(val);

answer[gid] = a[gid] + val;

}[/codebox]

I ever become errors.

How must I design the code to call such a kernel function?

A kernel function is not supposed to call other kernels. It appears that you want your kernel to call normal OpenCL functions, which can be declared as you have done above, but without the kernel qualifier. A side effect is that such functions, not being kernels, won’t be runnable from host code using clEnqueue__.

float add(float a, float b)

{

  return a+b;

}

__kernel void addVectors(__global float* many_a, __global float* many_b, __global float* result)

{

  size_t tid = get_global_id(0);

  float a=many_a[tid];

  float b=many_b[tid];

  result[tid] = add(a,b);

}

Have not tested this code-snipped so it may have minor syntax errors and such, but the general idea should be correct.

Yes, that works.

I’ve found out, that

float add(float a)

{

  return a+ 5;

}

and

float add(float a)

{

  return 5;

}

works, but

float add()

{

  return  5;

}

doesn’t.

A interesting side effect.

A kernel can absolutely call other kernels. Looking at this code:

[codebox]__kernel void

add five(__global float *a, __global float answer)

[/codebox]

This is not legal. No spaces are allowed within function names. Maybe this is your actual problem. Try:

[codebox]__kernel void

add_five(__global float *a, __global float answer)

[/codebox]

Another problem is kernels must both have a return type, which must be void. Saw both
__kernel __local float
five()

and

__kernel
five(__local float &val)