CUDA 流性能答案

【问题标题】：CUDA streams performanceCUDA 流性能
【发布时间】：2016-11-12 18:29:06
【问题描述】：

我目前正在通过计算两个向量之间的点积来学习 CUDA 流。成分是一个内核函数，它接受向量 x 和 y 并返回一个大小等于块数的向量 result，其中每个区块贡献了它自己减少的总和。

我还有一个宿主函数dot_gpu，它调用内核并将向量result缩减为最终的点积值。

同步版本就是这样做的：

// copy to device
copy_to_device<double>(x_h, x_d, n);
copy_to_device<double>(y_h, y_d, n);

// kernel           
double result = dot_gpu(x_d, y_d, n, blockNum, blockSize);

而异步是这样的：

double result[numChunks];
for (int i = 0; i < numChunks; i++) {
    int offset = i * chunkSize;

    // copy to device
    copy_to_device_async<double>(x_h+offset, x_d+offset, chunkSize, stream[i]);
    copy_to_device_async<double>(y_h+offset, y_d+offset, chunkSize, stream[i]);

    // kernel
    result[i] = dot_gpu(x_d+offset, y_d+offset, chunkSize, blockNum, blockSize, stream[i]);
}
for (int i = 0; i < numChunks; i++) {
    finalResult += result[i];
    cudaStreamDestroy(stream[i]);
}

我在使用流时性能越来越差，并试图调查原因。我尝试管道下载、内核调用和上传，但没有结果。

// accumulate the result of each block into a single value
double dot_gpu(const double *x, const double* y, int n, int blockNum, int blockSize, cudaStream_t stream=NULL)
{
double* result = malloc_device<double>(blockNum);
dot_gpu_kernel<<<blockNum, blockSize, blockSize * sizeof(double), stream>>>(x, y, result, n);

#if ASYNC
    double* r = malloc_host_pinned<double>(blockNum);
    copy_to_host_async<double>(result, r, blockNum, stream);

    CudaEvent copyResult;
    copyResult.record(stream);
    copyResult.wait();
#else
    double* r = malloc_host<double>(blockNum);
    copy_to_host<double>(result, r, blockNum);
#endif

double dotProduct = 0.0;
for (int i = 0; i < blockNum; i ++) {
    dotProduct += r[i];
}

cudaFree(result);
#if ASYNC
    cudaFreeHost(r);
#else
    free(r);
#endif

return dotProduct;
}

我的猜测是问题出在 dot_gpu() 函数内部，这些函数不仅调用内核。告诉我是否正确理解了以下流执行

foreach stream {
    cudaMemcpyAsync( device[stream], host[stream], ... stream );
    LaunchKernel<<<...stream>>>( ... );
    cudaMemcpyAsync( host[stream], device[stream], ... stream );
}

主机在没有被阻塞的情况下执行所有三个指令，因为 cudaMemcpyAsync 和内核立即返回（但是在 GPU 上它们将按顺序执行，因为它们被分配到同一个流）。所以主机继续下一个流（即使stream1知道它处于哪个阶段，但谁在乎......它在GPU上做他的工作，对吧？）并再次执行三个指令而不会被阻塞......等等等等。但是，我的代码在 dot_gpu() 函数内部的某个位置处理下一个流之前阻塞了主机。是不是因为我在分配和释放东西，以及将内核返回的数组减少到单个值？

【问题讨论】：

标签： cuda dot-product cuda-streams

【解决方案1】：

假设您的对象化 CUDA 接口执行函数和方法名称所暗示的那样，有以下三个原因导致后续调用 dot_gpu() 的工作可能不会重叠：

您的代码通过记录事件并等待它来显式阻止。
如果它没有阻塞 1.，您的代码将 block on the pinned host side allocation and deallocation，正如您所怀疑的那样。
如果您的代码没有阻塞 2.，则后续调用 dot_gpu() 的工作可能仍不会重叠，具体取决于计算能力。 Devices of compute capability 3.0 or lower do not reorder operations 即使它们被排入不同的流。

即使对于计算能力为 3.5 或更高的设备the number of streams whose operations can be reordered is limited by the CUDA_DEVICE_MAX_CONNECTIONS environment variable，它默认为 8，并且可以设置为最大 32 的值。

【讨论】：