CUDA inter-kernel communication between different streams

0
votes

Has anyone successfully run 2 different kernels in 2 different CUDA streams and gotten them to synchronize? Basically I want to have 1 kernel A send data to another concurrently running kernel B (in a different stream), then get results back. The reason: kernel A is running in 1 CUDA thread and I want a multiple GPU thread implementation for kernel B.

This is with high end GPUs (Fermi/Tesla), CUDA 4.2

Same GPU, different streams. So the data should be able to be communicated thru device memory, but how to sync them?

4
cudasync
You asking if a specific solution will work without stating your problem. In this particular instance you will find it better to state your problem so that the community can offer you solutions that are supported by the CUDA Programming Model. - Greg Smith

4 Answers

2
votes

The CUDA Programming Model only supports communication between threads in the same thread block (CUDA C Programming Guide at the end of section 2.2 Thread Hierarchy). This cannot be reliably implemented through the current CUDA API. If you try you may find partial success. However, this will fail on different OSes, different executions of your application, and this will be broken by future driver updates and new hardware (GK110 supports enhanced concurrency model).

0
votes

You will need to synchronize on the host. From the top of my head, calling cudaDeviceSynchronize for every stream in turn should do the trick but it may not be that easy.

0
votes
  • Your data must be in global memory
  • You need to get the data address on the host
  • You must send this data back to the second kernel

your code must be something similar to this:

*dataToExchange_h,*dataToExchange_d;
cudaMalloc((void**)dataToExchange, sizeof(data));

kernel1<<< M1,N1,0,stream1>>>(dataToExchange);
cudaStreamSynchronize(stream1);
kernel2<<< M2,N2,0,stream2>>>(dataToExchange);

But note that stream synchronization slow down the process, so you should avoid it as much as possible. You can also get stream synchronization through cuda events, it less obvious and does not give special advantage, but it's useful to know it ;-)

0
votes

If I correctly caught your question, you have two problems:

  1. Inter-Kernel data exchange
  2. Inter-Kernel synchronization

1) Inter-Kernel Data Exchange can be achieved through sharing data in global device memory.

2) As I know, there is no reliable facilities for inter-kernel synchronization provided by CUDA. And I'm unaware about any suitable trick that can be applied here.

CUDA C Programming Gide v7.5 tells us: "Applications manage the concurrent operations described above through streams. A stream is a sequence of commands (possibly issued by different host threads) that execute in order. Different streams, on the other hand, may execute their commands out of order with respect to one another or concurrently; this behavior is not guaranteed and should therefore not be relied upon for correctness (e.g., inter-kernel communication is undefined)."