该篇文章翻译参考自：https://leimao.github.io/blog/CUDA-Kernel-Execution-Overlap/
前提
在我的上一篇博客文章“CUDA Stream”中，我讨论了CUDA流如何帮助CUDA程序实现并发。在文章的结尾，我还提到除了内存传输和核函数执行重叠外，不同核函数之间的执行重叠也是允许的。然而，许多CUDA程序员想知道为什么他们以前没有遇到过核函数执行重叠。
在这篇博客文章中，我想讨论CUDA核函数执行重叠出现的条件以及为啥我们通常遇不到。
CUDA Kernel Execution Overlap
计算资源
如果有足够的计算资源来并行化多个CUDA内核执行，那么CUDA内核执行可以重叠。
在下面的示例中，通过将blocks_per_grid的值从小变大，我们可以看到不同CUDA流中的内核执行从完全并行化变为部分并行化，到最终几乎没有并行化。这是因为当分配给一个CUDA内核的计算资源变得更大时，额外的CUDA内核所需计算资源就会变得更少。
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#define CHECK_CUDA_ERROR(val) check((val), #val, __FILE__, __LINE__)
template <typename T>
void check(T err, const char* const func, const char* const file,
const int line)
{
if (err != cudaSuccess)
{
std::cerr << "CUDA Runtime Error at: " << file << ":" << line
<< std::endl;
std::cerr << cudaGetErrorString(err) << " " << func << std::endl;
std::exit(EXIT_FAILURE);
}
}
#define CHECK_LAST_CUDA_ERROR() checkLast(__FILE__, __LINE__)
void checkLast(const char* const file, const int line)
{
cudaError_t err{cudaGetLastError()};
if (err != cudaSuccess)
{
std::cerr << "CUDA Runtime Error at: " << file << ":" << line
<< std::endl;
std::cerr << cudaGetErrorString(err) << std::endl;
std::exit(EXIT_FAILURE);
}
}
__global__ void float_add_one(float* buffer, uint32_t n)
{
uint32_t const idx{blockDim.x * blockIdx.x + threadIdx.x};
uint32_t const stride{blockDim.x * gridDim.x};
for (uint32_t i{idx}; i < n; i += stride)
{
buffer[i] += 1.0F;
}
}
void launch_float_add_one(float* buffer, uint32_t n,
dim3 const& threads_per_block,
dim3 const& blocks_per_grid, cudaStream_t stream)
{
float_add_one<<<blocks_per_grid, threads_per_block, 0, stream>>>(buffer, n);
CHECK_LAST_CUDA_ERROR();
}
int main(int argc, char** argv)
{
size_t const buffer_size{1024 * 10240};
size_t const num_streams{5};
dim3 const threads_per_block{1024};
// Try different values for blocks_per_grid
// 1, 2, 4, 8, 16, 32, 1024, 2048
dim3 const blocks_per_grid{32};
std::vector<float*> d_buffers(num_streams);
std::vector<cudaStream_t> streams(num_streams);
for (auto& d_buffer : d_buffers)
{
CHECK_CUDA_ERROR(cudaMalloc(&d_buffer, buffer_size * sizeof(float)));
}
for (auto& stream : streams)
{
CHECK_CUDA_ERROR(cudaStreamCreate(&stream));
}
for (size_t i = 0; i < num_streams; ++i)
{
launch_float_add_one(d_buffers[i], buffer_size, threads_per_block,
blocks_per_grid, streams[i]);
}
for (auto& stream : streams)
{
CHECK_CUDA_ERROR(cudaStreamSynchronize(stream));
}
for (auto& d_buffer : d_buffers)
{
CHECK_CUDA_ERROR(cudaFree(d_buffer));
}
for (auto& stream : streams)
{
CHECK_CUDA_ERROR(cudaStreamDestroy(stream));
}
return 0;
}
$ nvcc overlap.cu -o overlap
$ ./overlap
我们观察到 blocks_per_grid = 1 时可以实现完全并行化。然而，我们也发现由于 GPU 没有被充分利用，完成所有内核的时间很长。
当我们设置 blocks_per_grid = 32 时，只有一部分内核执行是并行化的。然而，GPU 被充分利用，并且完成所有内核所花费的时间比 blocks_per_grid = 1 少得多。
与 blocks_per_grid = 32 相同，当我们将 blocks_per_grid = 5120 设置时，几乎没有内核执行并行化。然而，GPU仍然被充分利用，并且完成所有内核所花费的时间比 blocks_per_grid = 1 少得多。
Implicit Synchronization
即使有足够的计算资源，也可能不会发生内核执行重叠。这可能是由于主机线程发出的CUDA命令与其他不同流中的其他CUDA命令之间存在隐式同步。
在我看来，在单线程CUDA程序中很少发生这种情况，这是CUDA程序员通常编写CUDA程序的方式。然而，在多线程CUDA程序中肯定会发生这种情况。为了克服这种情况，自从CUDA 7以来，已经创建了一个“每个线程”默认流编译模式。用户只需在NVCC编译器构建标志中指定--default-stream per-thread即可，无需更改现有的CUDA程序以禁用隐式同步。要了解如何使用“每个线程”默认流简化CUDA并发性的更多详细信息，请阅读Mark Harris 的博客文章。
截至 CUDA 11.4，默认构建参数仍然是“legacy”。用户必须手动将其更改为“per-thread”，才能使用“每个线程”默认流。从 CUDA 11.4 NVCC 帮助文档：
--default-stream {legacy|null|per-thread} (-default-stream)
Specify the stream that CUDA commands from the compiled program will be sent
to by default.
legacy
The CUDA legacy stream (per context, implicitly synchronizes with
other streams).
per-thread
A normal CUDA stream (per thread, does not implicitly
synchronize with other streams).
'null' is a deprecated alias for 'legacy'.
Allowed values for this option: 'legacy','null','per-thread'.
Default value: 'legacy'.
结论
如果默认的CUDA流没有隐式同步，部分或无CUDA内核执行并行化通常表示GPU利用率高，并且完全的CUDA内核执行并行化通常表示GPU可能未被充分利用。
如果没有CUDA内核执行重叠是由于默认的CUDA流中的隐式同步，则应该考虑通过启用“per-thread”默认流来禁用它。