为什么不同流中的内核执行不是并行的？

Question

我刚刚在CUDA中学习了流技术，并尝试了它。然而，不希望的结果返回，即，流不是并行的。（在GPU Tesla M6，OS Red Hat Enterprise Linux 8上）

我有一个大小为（5,2048）的数据矩阵，还有一个处理矩阵的内核。

我的计划是分解“ nStreams = 4”扇区中的数据，并使用4个流来并行执行内核。

我的部分代码如下所示：

int rows = 5;
int cols = 2048;

int blockSize = 32;
int gridSize = (rows*cols) / blockSize;
dim3 block(blockSize);
dim3 grid(gridSize);

int nStreams = 4;    // preparation for streams
cudaStream_t *streams = (cudaStream_t *)malloc(nStreams * sizeof(cudaStream_t));
for(int ii=0;ii<nStreams;ii++){
    checkCudaErrors(cudaStreamCreate(&streams[ii]));
}

int streamSize = rows * cols / nStreams;
dim3 streamGrid = streamSize/blockSize;

for(int jj=0;jj<nStreams;jj++){
    int offset = jj * streamSize;
    Mykernel<<<streamGrid,block,0,streams[jj]>>>(&d_Data[offset],streamSize);
}    // d_Data is the matrix on gpu

Visual Profiler结果显示4个不同的流不是并行的。流13是第一个工作的流，而流16是最后一个工作的流。流13和流14之间有12.378us。每个内核执行大约持续5us。在上面的“运行时API”行中，它显示为“ cudaLaunch”。

你能给我一些建议吗？谢谢！

（我不知道如何在stackoverflow中上传图片，所以我只用文字描述结果。）

Answer 1

首先，不能保证在单独的流中启动的内容实际上将在GPU上并行执行。正如编程指南中指出的那样，使用多个流只会打开可能性，您不能依靠它实际发生的事情。由驾驶员决定。

除此之外，如果我没记错的话，您的Tesla M6有12个多处理器。这12个Maxwell多处理器中的每一个最多可容纳32个驻留块。这样，整个设备上驻留的最大块总数达到384。您正在启动320个块，每个块32个线程。仅此一项并不会留下太多的空间，您可能每个线程使用32个以上的寄存器，因此通过这些启动中的一个启动，GPU将会很满，这很可能是驱动程序选择不运行另一个内核的原因在平行下。

并行内核启动主要在您拥有时才有意义，例如，一堆执行不同任务的小内核可以在单独的多处理器上相邻运行。看来您的工作量可以轻松填满整个设备。您究竟希望通过并行运行多个内核来实现什么？为什么要使用这么小的模块？将整个程序作为具有更大块的大内核启动会更有意义吗？通常，您希望每个块至少有几个扭曲。参见，例如，更多问题：如何为CUDA内核选择网格和块尺寸？ 如果您使用共享内存，那么每个多处理器还需要至少两个块，否则您将无法在某些GPU上使用所有块（例如，每个多处理器提供96 KiB共享内存，但是每个GPU区块最多只能有48 KiB）…