人工智能系统性能工程 (Chinese Edition)

第11章. 核间 流水线处理、同步机制与CUDA流序内存分配

本作品已使用人工智能进行翻译。欢迎您提供反馈和意见：[e-mail]

迄今为止，我们主要探讨了单内核工具——cuda::pipeline 双缓冲、波束特化（加载/计算/存储波束）、持久内核以及带DSMEM/TMA的线程块集群——以保持单个内核在SM中的持续运行。本章将延续这些内核，展示如何通过CUDA流、事件及流顺序内存分配器实现跨内核与批次的流水线处理。简言之，第10章侧重于隐藏内核内部的延迟。本章则展示如何隐藏内核间以及GPU与主机间的延迟。

这种跨内核并行机制对于在实际工作负载中保持GPU所有引擎持续运转至关重要。要使现代GPU达到峰值利用率，必须让GPU的计算引擎与直接内存访问（DMA）引擎保持并行运行。

CUDA流为这种跨内核并行提供了基础。通过结合异步内存操作、精细同步机制以及CUDA图（本章简要介绍，下章详述），可构建高效管道以避免主机端阻塞。

利用CUDA流实现内核执行重叠

CUDA流是一系列 操作（内核启动、内存复制和内存分配）的序列，按发出顺序执行。如图11-1所示，设想从CPU向GPU启动5个内核，使用2个流。

图11-1. 从CPU向GPU上运行的两个流启动 五个内核

此时可见：ker_A 与ker_B 在流2上运行，而ker_1 、ker_2 及ker_3 在流1上运行。只要硬件资源允许，所有内核均可相互重叠——甚至跨CUDA流重叠。

CPU可在流异步执行内核操作期间继续处理任务（cpu_code_1 和cpu_code_2) ）。在两个CUDA流上启动这五个内核的代码如下所示：

#include <cstdio>
#include <cuda_runtime.h>

__global__ void ker_A()  { /* ... do some work ... */ }
__global__ void ker_B()  { /* ... do some work ... */ }

__global__ void ker_1()  { /* ... do some work ... */ }
__global__ void ker_2()  { /* ... do some work ... */ }
__global__ void ker_3()  { /* ... do some work ... */ }

int main() {
    // 1) Create two CUDA streams
    cudaStream_t stream1, stream2;
    cudaStreamCreateWithFlags(&stream1, cudaStreamNonBlocking);
    cudaStreamCreateWithFlags(&stream2, cudaStreamNonBlocking);

    // 2) Define your grid/block sizes
    dim3 grid(128);
    dim3 block(256);

    // 3) Launch ker_1 on stream1
    ker_1<<<grid, block, 0, stream1>>>();

    // 4) CPU code 1 runs ...