2、CUDA基础回顾：GPU执行模型、线程层次、内存层次、同步与异步的基本概念

在深入TensorRT的多流并发设计之前，必须扎实理解CUDA的底层执行模型。TensorRT的流（Stream）管理、异步推理、内存复用等高级特性，本质上都是对CUDA运行时（Runtime API）与驱动模型（Driver API）的封装与优化。如果对GPU的线程组织、内存层级和同步语义缺乏清晰认知，后续的并发调优将无从下手。本节将系统梳理CUDA编程中最核心的四个支柱：执行模型、线程层次、内存层次以及同步/异步机制，为后续多流并发设计奠定坚实的理论基础。

上图将CUDA的核心抽象归纳为四个相互关联的维度。执行模型决定了硬件如何调度计算，线程层次定义了软件的组织方式，内存层次划定了数据访问的路径与成本，而同步与异步机制则控制了并发执行的时序关系。TensorRT的多流并发设计，本质上就是在这四个维度上寻找最优的平衡点。下面逐一展开。

2.1 GPU执行模型：从SP到SM的硬件调度

核心原理：现代NVIDIA GPU（如Ampere、Hopper架构）由多个流多处理器（Streaming Multiprocessor, SM）组成。每个SM包含若干CUDA核心（即流处理器SP）、共享内存、寄存器文件和调度器。GPU执行的最小调度单位是线程束（Warp），通常包含32个线程。硬件以Warp为单位进行指令发射和执行，采用SIMT（单指令多线程）模型——即同一个Warp内的所有线程在同一时刻执行同一条指令，但可以操作不同的数据。

当一个CUDA Kernel启动时，GPU的工作调度器（GigaThread Engine）将线程块（Block）分配给空闲的SM。一个SM可以同时驻留多个Block，但受限于寄存器数量和共享内存容量。这种分配机制决定了占用率（Occupancy）——即SM上活跃Warp数量与最大可容纳Warp数量的比值。高占用率有助于隐藏内存访问延迟，但并非越高越好，过高的占用率可能导致寄存器溢出（Spill）到本地内存，反而降低性能。

2.2 线程层次：Grid-Block-Thread的三级组织

CUDA将线程组织为三层结构：网格（Grid）→ 线程块（Block）→ 线程（Thread）。Kernel启动时，程序员指定Grid和Block的维度（最多三维）。每个线程通过内置变量 blockIdx 和 threadIdx 定位自身身份。

• Grid：一次Kernel启动的全部线程集合。Grid中的Block数量可以远大于SM数量，硬件自动排队调度。
• Block：一组线程，块内线程可以通过共享内存协作，并通过 __syncthreads() 进行块内同步。Block是资源分配的基本单位。
• Thread：执行计算的最小单元。每个线程拥有独立的程序计数器和寄存器状态。

在TensorRT的推理引擎中，层（Layer）的Kernel通常由TensorRT自动生成，但理解线程层次有助于解释为什么某些层（如卷积、矩阵乘）在特定GPU上表现优异。例如，GEMM（通用矩阵乘）Kernel通常采用8x8或16x16的Block tile，以充分利用共享内存和寄存器重排。

2.3 内存层次：从全局到寄存器的访问代价

GPU内存层次的设计目标是隐藏延迟和最大化带宽。不同内存的访问速度差异巨大，理解这一点对优化TensorRT推理至关重要。

1. 全局内存（Global Memory）：位于显存（VRAM）中，容量最大（数GB至数十GB），但延迟高达数百个时钟周期。所有线程均可访问。访问模式必须满足合并访问（Coalesced Access）——即同一Warp的32个线程访问连续的内存地址，才能获得最佳带宽。
2. 共享内存（Shared Memory）：位于SM内部，容量小（通常几十KB至上百KB），延迟极低（约20-30周期）。块内线程共享，常用于数据复用和块内通信。TensorRT的卷积融合和Transformer的Attention计算大量依赖共享内存。
3. 寄存器（Register）：每个线程私有，访问速度最快（零延迟）。但数量有限（每个线程最多255个），超出后溢出到本地内存（Local Memory，实际位于全局内存中，性能极差）。
4. 常量/纹理内存：有专用缓存，适合广播访问或二维空间局部性访问。
5. L1/L2缓存：L1缓存与共享内存共享SM上的物理存储（可配置比例），L2缓存为所有SM共享。

2.4 同步与异步：CPU与GPU的协作哲学

同步（Synchronous）操作会阻塞CPU直到GPU完成指定任务。例如 cudaMemcpy（同步版本）会等待数据拷贝完成才返回。异步（Asynchronous）操作则立即返回控制权给CPU，GPU在后台执行。异步操作是CUDA并发编程的基石。

CUDA Stream（流）是管理异步操作的核心抽象。一个Stream是一个按顺序执行的命令序列（Kernel启动、内存拷贝、事件记录等）。不同Stream之间的操作可以并发执行（如果硬件资源允许）。TensorRT的多流并发正是基于这一机制：将不同的推理请求分配到不同的Stream中，实现GPU资源的时分复用和吞吐量提升。

同步原语包括：

• cudaDeviceSynchronize()：阻塞CPU直到所有Stream上的所有操作完成。
• cudaStreamSynchronize(stream)：阻塞直到指定Stream完成。
• cudaEventRecord(event, stream)：在Stream中记录一个事件，用于时间测量或跨流同步。
• cudaStreamWaitEvent(stream, event)：让一个Stream等待另一个Stream的事件，实现跨流依赖。

2.5 案例拆解：一个典型的异步推理流水线

假设我们要设计一个高吞吐的推理服务，输入为一批图像，输出为分类结果。初始方案是同步推理：每张图像依次执行“拷贝到GPU → 推理 → 拷贝回CPU”，耗时约10ms/张，GPU利用率仅40%。

优化目标：将吞吐量提升至原来的2倍以上，同时保持延迟在可接受范围（<20ms）。

分析过程：通过Nsight Systems分析发现，GPU大部分时间处于空闲状态，等待CPU完成图像预处理（如Resize、Normalize）。CPU与GPU的工作完全串行。

操作步骤：

1. 创建两个CUDA Stream：stream1 和 stream2。
2. 将图像预处理（Resize、Normalize）也写成CUDA Kernel，在GPU上执行。
3. 使用 cudaMemcpyAsync 将原始图像异步拷贝到GPU固定内存（Pinned Memory）。
4. 在 stream1 中依次提交：预处理Kernel → TensorRT推理 → 结果拷贝回CPU（异步）。
5. 在 stream2 中提交另一批图像的处理，与 stream1 重叠执行。
6. 使用 cudaEvent 确保 stream2 的预处理不会覆盖 stream1 正在使用的输入缓冲区。

结果对比：优化后，GPU利用率提升至85%，吞吐量从100张/秒提升至220张/秒，平均延迟从10ms增加到14ms（因排队等待），但满足<20ms的约束。

风险与代价：多流并发增加了代码复杂度，需要仔细管理缓冲区生命周期。若两个Stream同时访问同一块全局内存，可能产生竞争条件。此外，并非所有GPU都支持充分的并发执行——消费级GPU（如RTX 30系列）的并发能力弱于数据中心GPU（如A100）。

复盘经验：多流并发的核心是让GPU始终有活干。在TensorRT中，可以通过 setStream() 方法将执行上下文绑定到特定Stream，并利用 cudaStreamAddCallback 或事件来通知CPU推理完成，从而避免忙等待。

2.6 常见误区与总结复盘