Kernel call from within another kernel


I have code that calls functions from within functions. So I want to know if I want to keep the same structure, is it possible that from within OpenCL kernel (as a function), I try to call another kernel (like function) and dont go back to Host program so that I dont have to make new buffers and copy data back and forth from kernel buffers and host program memories ?

I’m not sure I understand the question. In OpenCL calling functions inside kernels is allowed as long as the call is not recursive. Kernels can also call other kernel functions and the behavior is the same as calling any other (non-kernel) function.

You don’t have to go back to the host or copy memory to/from the host. memory stays on the device unless you ask to copy it, and the work queue is a queue anyway: returning from the kernel is not a synchronous call back to the host it’s just a chance for the device to call the next kernel in the queue.

Calling kernels from other kernels is of limited use because you (should) only use kernels to separate global data boundaries in the first place. i.e. if you are able to share code between the kernels you’d probably combine them anyway.

Thanks for all valuable comments. But would I have to call thread awakening function i.e.,
int x = get_global_id(0);
again inside every kernel in kernel calling hierarchy ?

It doesn’t awaken threads, it just returns a unique job identifier. I’m not sure what you even mean by that?

This can be called from functions your kernel invokes, including other kernels. They are treated (mostly) the same way as any other function.