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

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 个驻留块。这使得驻留在整个设备上的最大块总数达到 384。您将启动 320 个块，每个块有 32 个线程。仅此一项并没有留下那么多空间，而且每个线程可能使用超过 32 个寄存器，因此 GPU 将在其中一次启动时非常满，这很可能是驱动程序选择不运行另一个内核的原因并行。

并行内核启动主要是在您拥有的情况下才有意义，例如，一堆小内核执行不同的工作，这些内核可以在不同的多处理器上彼此相邻运行。您的工作量似乎可以轻松填满整个设备。您希望通过并行运行多个内核来实现什么？你为什么要处理这么小的块？将整个事物作为一个具有更大块的大内核来启动不是更有意义吗？通常，您希望每个块至少有几个经线。例如，请参阅此问题了解更多信息：How do I choose grid and block dimensions for CUDA kernels? 如果您使用共享内存，您还需要每个多处理器至少两个块，否则您甚至无法在某些 GPU 上使用所有块（其中，例如，为每个多处理器提供 96 KiB 共享内存，但每个块最多只能有 48 KiB）...

Answer 2

要添加到现有答案（这是完全正确的），请考虑以下您在问题中发布的代码的简单完整版本：

__global__
void Mykernel(float* data, int size)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    for(; tid < size; tid+= blockDim.x * gridDim.x) data[tid] = 54321.f;
}

int main()
{
    int rows = 2048;
    int cols = 2048;

    int blockSize = 32;
    dim3 block(blockSize);

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

    float* d_Data;
    cudaMalloc(&d_Data, sizeof(float) * rows * cols);
    int streamSize = rows * cols / nStreams;
    dim3 streamGrid = dim3(4);

    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


    cudaDeviceSynchronize();
    cudaDeviceReset();
}

请注意两个区别——通过将rows 设置为 2048，减少了每个内核启动的块数，并且增加了每个线程的总计算量。内核本身包含一个网格步长循环，它允许每个线程处理多个输入，确保无论启动多少块/线程，都能处理整个输入数据集。

在与您的设备类似的 Maxwell GPU 上进行分析表明：

即内核确实重叠。现在让我们将问题大小减小到问题中指定的大小（行 = 5）：

内核不再重叠。为什么？因为驱动程序和设备延迟足够高，并且每个内核的执行时间足够短，以至于没有时间发生执行重叠，即使设备资源允许它。因此，除了另一个答案中描述的资源要求限制之外，计算量必须足够大，以抵消与在流中调度内核启动相关的固定延迟。

最后，我建议设置基于流的并发执行方案的正确方法应该如下所示：

int blockSize = 32;
dim3 block(blockSize);
int blocksperSM, SMperGPU = 13; // GPU specific
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&blocksperSM, Mykernel, blockSize, 0); // kernel specific
dim3 streamGrid = blocksperSM * (SMperGPU / nStreams); // assume SMperGPU >> nstreams   

这里的想法是，可用 SM 的数量（大致）在流中平均分配，并且通过占用 API 为内核获取所选块大小的最大占用每个 SM 的块数。

此配置文件如下：

在rows = 2048 的情况下，通过将内核的资源需求与 GPU 的容量正确匹配，从而产生重叠和较短的执行时间。