Say, I have two different kernels, kernel 0 and kernel 1, running at the same time on the same GPU. Let's say, each kernel is exactly 64 threads/1 wave. At some point, kernel 1 wants to send any kind of signal - it can even be a 1 bit set to 1 - to kernel 0. Is this somehow possible without SVM/OpenCL 2.0?
I can send a signal to the CPU with s_sendmsg and deal with an IRQ on the CPU side, but I don't see a way to send a signal within the GPU, or even send a signal from the CPU to the GPU while a kernel is running, e.g. by setting a bit to 1 in a memory buffer. (If the waves belong to the same kernel launched with the same clEnqueueNDRangeKernel call, I obviously can use GDS and many other things.)
So, is wave to wave within the same GPU when waves belong to different kernels or CPU to GPU signaling possible while the kernels are running? If so, how do I do that?
As per OpenCL standard, above scenarios are not possible without SVM atomics support.