CUDA101 - 05 异步并行-底层控制与系统优化

2026/04/12

  • 线程束特化（Warp Specialization）
    • 核心思想是：在同一个 thread block 内，让不同的 warp 执行不同类型的任务，而不是所有线程执行完全相同的指令。典型做法是将 warp 分为"生产者"和"消费者"，例如一个 warp 负责从全局内存加载数据到共享内存（load），另一个 warp 负责使用这些数据进行计算（compute），从而实现访存与计算的重叠。
    • 实现方式本质上就是基于 warp id 的分支控制。通过计算 warp_id = threadIdx.x / 32（或更通用地基于线性线程 id），在 kernel 内用条件语句区分不同 warp 的职责，例如 warp_id == 0 的 warp 执行访存，其他 warp 执行计算。这种分支不会导致性能问题，因为 warp 内的线程具有相同的 warp_id，不会发生 warp 内分歧，分支只发生在 warp 之间。
    • 关键问题在于如何避免数据冲突。warp specialization 通常配合双缓冲（double buffering）使用，将共享内存划分为两个 buffer。某一时刻，producer warp 向 buffer A 写入数据，同时 consumer warp 从 buffer B 读取并计算；下一轮交换角色。这样读写发生在不同 buffer 上，从根本上避免了读写冲突。同时需要配合同步机制保证数据可见性：粗粒度可以使用 __syncthreads()，但会阻塞；更高效的方式是使用 pipeline（如 cuda::pipeline 或 async memcpy 的 commit/wait 机制），实现细粒度的 producer-consumer 同步。
    • 整体执行可以看作一个软件流水线：在第 t 步加载 tile t+1 的同时，计算 tile t，从而把原本串行的"load → compute"转化为重叠执行。这一机制的本质是利用 GPU 中访存单元和计算单元的独立性，将时间维度上的空闲转化为并行。
  • CUDA pipeline
    • CUDA pipeline 是一种用于实现存算重叠的编程抽象，本质是把"数据拷贝（访存）"与"计算"组织成一个多阶段的软件流水线，从而隐藏全局内存访问的高延迟。它是对前面 warp specialization 和双缓冲思路的系统化封装。
    • 从问题出发：在 warptiling 优化后，kernel 仍然呈现出明显的 latency-bound 特征，表现为 compute 与 memory 利用率都不高，大量时间消耗在 Long Scoreboard Stall（等待全局内存）。根本原因是执行流程仍然是串行的 load → compute，没有时间上的重叠。
    • pipeline 的核心思想是将执行拆成多个 stage（例如 stage=2 对应双缓冲），在时间上形成流水：当前阶段进行计算的同时，下一阶段的数据已经在异步加载。这样原本串行的访存与计算被转化为并行执行，从而实现延迟隐藏。
    • 在抽象上，pipeline 将线程分为两个逻辑角色：producer 负责发起数据拷贝（GMEM → SMEM），consumer 负责对已准备好的数据进行计算。这种 producer-consumer 关系不再必须通过 warp specialization 显式绑定，而是通过 pipeline 机制在逻辑上表达。
    • pipeline 的执行流程可以概括为：acquire → memcpy_async（submit）→ commit → wait → compute → release。其含义是：先申请一个可用的 stage buffer，然后发起异步拷贝并提交，consumer 在需要使用数据前进行 wait，同步保证数据已就绪，随后执行计算，最后释放该 stage 供后续迭代复用。整个过程中，memcpy_async + compute 是可以并行的，而 acquire / commit / wait / release 属于协作同步点。
    • 在实现上，pipeline 通常与 __pipeline_memcpy_async（或底层 cp.async）配合使用，该接口可以将数据直接从全局内存异步搬运到共享内存，且要求拷贝粒度为 4/8/16 字节。相比传统同步加载（ld.global → st.shared），这种方式避免了线程阻塞等待，从而显著降低 Long Scoreboard Stall。
    • 与之前的优化手段对应关系如下：双缓冲对应 pipeline 的多 stage；warp specialization 中的 load/compute 分工对应 producer/consumer；__syncthreads() 这样的全局同步被更细粒度的 wait/commit 替代；手动管理 buffer 切换被 pipeline 的 acquire/release 机制封装。pipeline 的优势在于将这些模式统一为一个可组合的异步执行框架，既支持单 warp 内部的流水，也支持更复杂的线程协作。
    • 整体执行可以理解为：在第 t 次迭代中计算 tile t 的同时，已经在异步加载 tile t+1，下一轮直接使用已就绪的数据继续计算，从而将访存延迟"折叠"进计算时间中。
  • pipeline的发起与提交
    • __pipeline_memcpy_async 与 __pipeline_commit 构成 pipeline 的"发起-提交"机制。前者用于声明一次从 GMEM 到 SMEM 的异步拷贝请求，将该请求加入当前 pipeline stage 的待执行队列；后者用于将当前 stage 中累计的所有拷贝请求一次性提交给硬件执行。两者对应 pipeline 抽象中的 submit 与 commit 阶段。这样设计的目的是支持批量发射访存操作，并明确划分不同 stage 的边界，从而构建多级流水线。若没有 commit，memcpy_async 的请求不会真正执行。整体流程为：memcpy_async（登记）→ commit（发射）→ wait（等待完成）→ compute（使用数据）。
  • 多流（multi-stream）
    • 多流是在 CUDA 中实现多任务并行的机制，本质是通过多个独立的执行队列（stream）让多个 kernel 在 GPU 上重叠执行，从而进一步隐藏延迟、提高资源利用率。一个 stream 可以看作一个 FIFO 队列，同一 stream 内的操作严格顺序执行，不同 stream 之间在资源允许的情况下可以并行执行。
    • 默认情况下，所有 kernel 都被提交到默认流（stream 0），因此即使代码中没有显式使用 stream，也能正常运行，但此时所有操作是串行的。更重要的是，默认流具有隐式同步语义：默认流中的操作会与其他显式流发生同步，因此"默认流 + 显式流"通常仍然是顺序执行的，无法实现并行。
    • 要实现真正的多流并行，需要显式创建多个 stream，并将不同 kernel 分配到不同 stream 中执行。同时，为避免默认流带来的隐式同步，需要使用 cudaStreamCreateWithFlags(..., cudaStreamNonBlocking) 创建非阻塞流，使其可以与默认流以及其他流独立并行。
    • 从优化层级上看，多流解决的是"多个任务之间"的并行问题，而 pipeline / async memcpy 解决的是"单个 kernel 内部"的存算重叠问题。前者属于粗粒度并行（kernel 级），后者属于细粒度并行（指令 / 数据级）。多流的核心作用是在单个 kernel 无法完全占满 GPU 资源或存在等待时，通过并发执行多个 kernel 来提高整体吞吐。整体目标仍然是延迟隐藏：当一个 kernel 在等待访存或同步时，其他 kernel 可以继续执行，从而避免 GPU 资源空闲。
  • CPU与GPU间数据搬运的异步执行
    • __pipeline_memcpy_async 与 cudaMemcpyAsync / cudaFreeAsync 属于两类不同层级的异步机制。前者作用于 kernel 内部，用于实现 GMEM→SMEM 的异步拷贝，配合 pipeline（commit / wait）实现计算与访存的重叠，解决的是单个 kernel 内的延迟隐藏问题。后者作用于 kernel 外部（CPU↔GPU），基于 stream 调度，将数据传输或内存释放操作加入执行队列，使其与 kernel 执行重叠，解决的是任务级（kernel 与 memcpy、kernel 与 kernel）之间的延迟隐藏问题。
    • cudaMemcpyAsync 的本质是将数据拷贝操作异步提交到某个 stream 中，使 CPU 不阻塞，并允许 memcpy 与 kernel 在不同 stream 上并行执行；cudaFreeAsync 则将内存释放延迟到对应 stream 中所有相关操作完成之后再执行，从而避免全局同步。相比之下，__pipeline_memcpy_async 只在 GPU 内部生效，不涉及 CPU 调度。
    • 整体可以分为三层并行与异步：kernel 内部通过 pipeline 实现访存与计算重叠（细粒度），kernel 之间通过 multi-stream 实现并行执行（中粒度），CPU 与 GPU 之间通过 cudaMemcpyAsync 等接口实现数据传输与计算重叠（粗粒度）。三者共同构成完整的延迟隐藏体系。
  • cudaMallocAsync与cudaMallocHost
    • cudaMemcpyAsync 与 pinned memory 相关，决定的是 CPU↔GPU 数据传输能否真正异步：只有当 host 内存是 pinned（cudaMallocHost 分配）时，GPU 才能通过 DMA 直接访问，从而实现真正的异步拷贝与传输-计算重叠；若使用 pageable memory（malloc），则会在内部引入同步的 staging copy，cudaMemcpyAsync 会退化为"伪异步"。
    • cudaMallocAsync 与此无关，它作用于 GPU 端内存管理，用于在 stream 中异步进行 device memory 的分配与释放，避免传统 cudaMalloc/cudaFree 带来的全局同步开销。
    • 两者分别解决不同层级的问题：前者是数据传输的异步化，后者是 GPU 内存管理的异步化，不能互相替代。
  • Unified Memory（Managed Memory）
    • Unified Memory（Managed Memory）通过 cudaMallocManaged 分配，使 CPU 和 GPU 共享同一虚拟地址空间，程序员无需显式调用 cudaMemcpy，数据在两者之间的迁移由系统自动完成。其底层机制是按页（page）迁移：内存被划分为固定大小的页（如 4KB），当某一侧（CPU 或 GPU）访问数据而该页不在本地时，会触发 page fault，由驱动将该页迁移到当前访问方。
    • Page fault 的本质是"访问缺页触发迁移"，在 GPU 场景中表现为：GPU 访问某页数据时若该页仍在 CPU 侧，则产生 page fault，驱动发起 DMA 将数据搬运到 GPU，随后继续执行。该过程会引入较高延迟，并可能中断 kernel 执行；若访问模式频繁在 CPU 与 GPU 间切换，还可能导致页在两侧反复迁移（thrashing）。
    • Unified Memory 简化了编程模型，但性能可控性较差。可通过 cudaMemPrefetchAsync 预取数据、cudaMemAdvise 提供访问提示来减少 page fault。整体上，它与 pageable/pinned memory 的区别在于：pageable/pinned 依赖显式 memcpy，而 unified memory 依赖隐式按页迁移。
  • CUDA Graph
    • CUDA Graph 将一组 GPU 操作（如 kernel、memcpy）组织成一个有向无环图（DAG），用节点表示任务、边表示依赖关系。其核心思想是将原本通过 stream 逐条提交的操作序列"录制"为一个整体执行单元（通过显式构建或 stream capture），再通过一次 cudaGraphLaunch 提交执行，从而显著减少 CPU 端的 kernel launch 和调度开销。
    • 执行流程包括：构建/捕获 graph → cudaGraphInstantiate 实例化 → 多次 cudaGraphLaunch 复用执行。相比传统方式，CUDA Graph 将多次 API 调用合并为一次提交，提升调度效率，特别适用于结构固定、重复执行、小 kernel 较多的场景。其优化层级属于"CPU 调度层"，与 kernel 内的 pipeline、kernel 间的 stream 并行互补。
  • 多GPU通信 & NCCL
    • 多 GPU 通信的核心问题是在不同 GPU 之间高效交换数据，常见路径包括经由主机内存（GPU↔CPU↔GPU）和设备直连（如 NVLink / PCIe P2P）。相较传统 cudaMemcpy 的点对点拷贝，通信模式通常分为 collective（集体通信）操作，如 broadcast、reduce、all-reduce、all-gather 等，用于多 GPU 协同计算中的数据同步与聚合。
    • NCCL（NVIDIA Collective Communications Library）提供了针对多 GPU / 多节点的高性能通信实现，自动选择最优拓扑（如 ring、tree）和传输路径（NVLink、PCIe、InfiniBand），并与 CUDA stream 深度集成，支持与计算 overlap。其典型接口如 ncclAllReduce、ncclBroadcast 等，通常在同一 stream 中发起，实现通信与 kernel 的流水化执行。
    • 性能关键在于带宽利用与通信-计算重叠：通过分块（chunking）+ pipeline，使数据在传输过程中逐步参与计算；通过多 stream 或异步调度避免全局同步；通过拓扑感知（如 ring all-reduce）最大化链路带宽。NCCL 屏蔽了底层细节，将多 GPU 通信抽象为高效的 collective 原语，是分布式训练中的核心组件。