是否有可能同时在GPU上运行2个内核，并相互通信?

Is it possible to have 2 kernels running on GPU at same time, and communicating with each other?

本文关键字：内核 互通 通信 2个 运行 有可能 GPU 是否      更新时间：2023-10-16

如果我有两个相关的程序，我是否有可能让一个内核在一个SM上运行，另一个内核同时在其他SM(或SMs)上运行?此外，我需要他们能够通过全局内存相互通信。这可能吗?我可以使用cuda流来实现这一点吗?

理论上是可能的。

我认为这很容易出现问题，因为，与内核内全局同步不同，它依赖于两个内核能够达到各自的同步点。这通常意味着非常小的内核，您可以保证无论块的启动顺序如何，都将到达同步点。

但一个简单的证明可能的情况如下:

#include <stdio.h>
#include <unistd.h>
__device__ volatile unsigned int sem = 0;
__global__ void kernel1(){
  while(sem < 1);
  printf("kernel1 received signal, sending return signaln");
  sem = 2;
  __threadfence();
}
__global__ void kernel2(){
  printf("kernel2 sending signaln");
  sem = 1;
  __threadfence();
  while(sem<2);
  printf("kernel2 received signaln");
}

int main(){
  cudaStream_t stream1, stream2;
  cudaStreamCreate(&stream1);
  cudaStreamCreate(&stream2);
  printf("Launching kernel 1n");
  kernel1<<<1,1,0,stream1>>>();
  sleep(2);
  printf("Launching kernel 2n");
  kernel2<<<1,1,0,stream2>>>();
  cudaDeviceSynchronize();
  return 0;
}

对于cc2.0设备，这至少需要编译。

当运行时，它应该给出这样的输出:

Launching kernel 1
Launching kernel 2
kernel2 sending signal
kernel1 received signal, sending return signal
kernel2 received signal

另一方面，如果您从内核启动中删除stream1和stream2标识符，则程序将挂起(因为内核都启动到相同的流中，因此被序列化)

再次强调，我不认为这是一个好的设计，但在某些情况下是可能的。

如果你使用的是cc3.5或更新的GPU，你可能想要研究使用动态并行(从一个内核启动另一个内核)。

相关文章：