CUDA 中的异步流不会提高性能

Question

我正在尝试通过使用多个流来加速以下 CUDA 代码。

#define N (4096 * 4096)
#define blockDimX  16
#define blockDimY  16

float domain1 [N];
float domain2 [N];

__global__ updateDomain1_kernel(const int dimX, const int dimY) {
    // update mechanism here for domain1
    // ...
}

__global__ updateDomain2_kernel(const int dimX, const int dimY) {
    // update mechanism here for domain2, which is nearly the same
    // ...
}

__global__ addDomainsTogether_kernel(float* domainOut, 
                                     const int dimX, 
                                     const int dimY) 
{
    // add domain1 and domain2 together and fill domainOut
}

void updateDomains(float* domainOut) {
    dim3 blocks((dimX + blockDimX - 1) / blockDimX , (dimY + blockDimY- 1) / blockDimY);
    dim3 threads(blockDimX, blockDimY);

    updateDomain1_kernel<<<blocks, threads>>> (dimX, dimY);
    updateDomain2_kernel<<<blocks, threads>>> (dimX, dimY);
    addDomainsTogether_kernel<<<block, threads>>> (domainOut_gpu, dimX, dimY);
    cudaMemcpy(domainOut, domainOut_gpu, N * sizeof(float), cudaMemcpyDeviceToHost);
}

精确的实现并不重要；重要的是更新各自的域是两个完全独立的操作，之后都在第三个内核调用中使用。因此，我认为尝试通过将每个更新内核放在自己的流中来加速它是一个好主意，我想同时运行它。所以我把它改成了下面这样：

void updateDomains(float* domainOut) {
    dim3 blocks((dimX + blockDimX - 1) / blockDimX , (dimY + blockDimY- 1) / blockDimY);
    dim3 threads(blockDimX, blockDimY);

    cudaStream_t stream0, stream1;
    cudaStreamCreate(&stream0);
    cudaStreamCreate(&stream1);

    updateDomain1_kernel<<<blocks, threads, 0, stream0>>> (dimX, dimY);
    updateDomain2_kernel<<<blocks, threads, 0, stream1>>> (dimX, dimY);
    cudaDeviceSynchronize();

    addDomainsTogether_kernel<<<block, threads>>> (domainOut_gpu, dimX, dimY);
    cudaMemcpy(domainOut, domainOut_gpu, N * sizeof(float), cudaMemcpyDeviceToHost);

    cudaStreamDestroy(stream0);
    cudaStreamDestroy(stream1);
}

我推测发现了性能速度的差异，但绝对没有明显的差异。因此，考虑到编译器可能是第一次通过自动调度更新调用而变得聪明，我认为以下应该会降低性能：

void updateDomains(float* domainOut) {
    dim3 blocks((dimX + blockDimX - 1) / blockDimX , (dimY + blockDimY- 1) / blockDimY);
    dim3 threads(blockDimX, blockDimY);

    cudaStream_t stream0;
    cudaStreamCreate(&stream0);

    updateDomain1_kernel<<<blocks, threads, 0, stream0>>> (dimX, dimY);
    updateDomain2_kernel<<<blocks, threads, 0, stream0>>> (dimX, dimY);

    addDomainsTogether_kernel<<<block, threads0, stream0>>> (domainOut_gpu, dimX, dimY);
    cudaMemcpy(domainOut, domainOut_gpu, N * sizeof(float), cudaMemcpyDeviceToHost);

    cudaStreamDestroy(stream0);
}

但是，性能速度几乎没有任何差异。如果有的话，最后一个似乎最快。这让我觉得我不了解 CUDA 流。有人能告诉我如何加速这段代码吗？

Answer 1

如果您没有已经使用了所有可用的内核，增加的并行性只会增加您的计算吞吐量。如果您已经拥有足够的并行度，那么它只会增加同步开销。