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

Question

我刚刚学习了 CUDA 中的流技术，并且尝试了一下。然而，不希望的结果返回，即流不是平行的。 (在 GPU Tesla M6 上，操作系统 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。在上面的“Runtime API”行中，它表示“cudaLaunch”。

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

(我不会stackoverflow怎么上传图片，所以只用文字描述一下结果。)

Answer 1

首先，无法保证在单独的流中启动的内容实际上会在 GPU 上并行执行。作为pointed out in the programming guide ，使用多个流只是打开了可能性，你不能依赖它实际发生。由司机决定。

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

并行内核启动主要在您拥有时才有意义，例如，一堆做不同事情的小内核可以在单独的多处理器上彼此相邻运行。看起来您的工作量很容易填满整个设备。您究竟希望通过并行运行多个内核来实现什么？你为什么要用这么小的 block ？将整个东西作为一个具有更大块的大内核来启动不是更有意义吗？通常，您希望每个 block 至少有几个扭曲。有关更多信息，请参见例如这个问题:How do I choose grid and block dimensions for CUDA kernels?如果您使用共享内存，您还需要每个多处理器至少有两个 block ，否则您甚至无法在某些 GPU 上使用所有 block (例如，每个多处理器提供 96 KiB 共享内存，但每个 block 最多只能有 48 KiB)...