如果我有两个相关的程序,我是否有可能让一个内核在一个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,你可能想要研究使用动态并行(从一个内核启动另一个内核)。