124 lines
5.6 KiB
Markdown
124 lines
5.6 KiB
Markdown
|
|
---
|
|||
|
|
tags:
|
|||
|
|
date created: 星期四, 十一月 20日 2025, 9:50:03 晚上
|
|||
|
|
date modified: 星期四, 十一月 20日 2025, 9:50:24 晚上
|
|||
|
|
---
|
|||
|
|
|
|||
|
|
# 2.2.2 异步流水线与计算通信重叠 (Asynchronous Pipelining & Compute-Copy Overlap)
|
|||
|
|
|
|||
|
|
**覆盖范围**:定义如何利用智铠 GPU 的独立 Copy Engine (DMA 引擎) 与 Compute Engine (计算引擎) 的并行能力,通过 **CUDA Streams** 实现“传输 - 计算 - 传输”的三级流水线并行,从而掩盖 PCIe 总线的物理延迟。
|
|||
|
|
|
|||
|
|
#### 一、 约束输入与对齐
|
|||
|
|
|
|||
|
|
1. **硬件能力**:Iluvatar MR-V100 通常具备独立的 Copy Engine(用于 H2D/D2H)和 Compute Engine。这意味着 **数据拷贝** 和 **Kernel 执行** 在硬件上是物理隔离的,可以同时进行。
|
|||
|
|
2. **API 约束**:必须使用 **Async** 系列 API (如 `cudaMemcpyAsync`) 配合 **Non-Default Stream** 才能触发重叠。
|
|||
|
|
3. **业务逻辑**:雷达信号处理通常是流式的:`接收(H2D) -> 处理(Kernel) -> 输出(D2H)`。
|
|||
|
|
|
|||
|
|
#### 二、 权衡分析与选项呈现 (Trade-off Matrix)
|
|||
|
|
|
|||
|
|
我们主要在**流的设计模式**上进行权衡:
|
|||
|
|
|
|||
|
|
| 选项 | A. 单流串行 (Serial Stream) | B. 多流乒乓/多缓冲 (Multi-Stream Ping-Pong) **(推荐)** | C. 细粒度多流 (Hyper-Q) |
|
|||
|
|
| :--- | :--- | :--- | :--- |
|
|||
|
|
| **机制** | 1 个流。H2D -\> Kernel -\> D2H 顺序执行。 | 2-3 个流。Stream A 做计算时,Stream B 做 H2D 拷贝。 | N 个流(N \>\> 3)。将任务切分为极小片。 |
|
|||
|
|
| **PCIe 利用率** | **低**。总线在 Kernel 计算期间闲置。 | **高**。总线和计算单元始终处于忙碌状态。 | **极高**,但调度开销大。 |
|
|||
|
|
| **延迟掩盖** | 无掩盖。总耗时 = T(copy) + T(compute)。 | **完全掩盖**。理想情况下总耗时 = max(T(copy), T(compute))。 | 同上,但可能引入调度抖动。 |
|
|||
|
|
| **实现复杂度** | 低。 | 中。需要管理多个 Buffer 的状态 (Ping-Pong)。 | 高。 |
|
|||
|
|
| **适用性** | 调试模式。 | **雷达实时处理标准范式。** | 超大规模并发任务。 |
|
|||
|
|
|
|||
|
|
#### 三、 基线确立与实施规范
|
|||
|
|
|
|||
|
|
为了最大化吞吐量,我们确立 **B. 多流乒乓 (Multi-Stream Ping-Pong)** 为设计基线。
|
|||
|
|
|
|||
|
|
##### 1\. 流水线架构基线:三级流水线 + 双流 (Double Buffering)
|
|||
|
|
|
|||
|
|
- **核心逻辑**:创建 **2 个 CUDA Stream** (Stream 0, Stream 1) 和 **2 组页锁定内存 Buffer** (Buffer A, Buffer B)。
|
|||
|
|
- **调度策略**:
|
|||
|
|
- **时刻 T0**:Stream 0 开始传输 Buffer A (H2D)。
|
|||
|
|
- **时刻 T1**:
|
|||
|
|
- Stream 0 开始处理 Buffer A (Kernel)。
|
|||
|
|
- **同时**,Stream 1 开始传输 Buffer B (H2D) —— **此处发生了 Copy 与 Compute 的重叠**。
|
|||
|
|
- **时刻 T2**:
|
|||
|
|
- Stream 0 开始回传 Buffer A 结果 (D2H)。
|
|||
|
|
- Stream 1 开始处理 Buffer B (Kernel)。
|
|||
|
|
|
|||
|
|
##### 2\. 关键 API 实施规范
|
|||
|
|
|
|||
|
|
- **流创建**:
|
|||
|
|
|
|||
|
|
```cpp
|
|||
|
|
cudaStream_t streams[2];
|
|||
|
|
for(int i=0; i<2; i++) cudaStreamCreateWithFlags(&streams[i], cudaStreamNonBlocking);
|
|||
|
|
```
|
|||
|
|
|
|||
|
|
- *注意*:必须使用 `cudaStreamNonBlocking`,防止与默认流(Default Stream)发生隐式同步,导致流水线断流。
|
|||
|
|
- **异步传输**:
|
|||
|
|
|
|||
|
|
```cpp
|
|||
|
|
// 必须使用 Async 版本,且指定 stream
|
|||
|
|
cudaMemcpyAsync(d_ptr, h_ptr, size, cudaMemcpyHostToDevice, streams[i]);
|
|||
|
|
```
|
|||
|
|
|
|||
|
|
- **同步策略**:
|
|||
|
|
- **严禁**使用 `cudaDeviceSynchronize()`(全卡同步)。
|
|||
|
|
- **推荐**使用 `cudaStreamSynchronize(streams[i])` 或 `cudaEventRecord/Synchronize` 来精细控制单个 Buffer 的生命周期,确保 Host 端在复用 Buffer 前,GPU 已经操作完毕。
|
|||
|
|
|
|||
|
|
##### 3\. 缓冲区管理状态机
|
|||
|
|
|
|||
|
|
为了配合 `01_数据接收模块` 的 `MemoryPool`,我们需要一个简单的状态机来管理 Buffer 在 Host 和 Device 之间的流转:
|
|||
|
|
|
|||
|
|
- `HOST_OWNED` (I/O 线程填充数据)
|
|||
|
|
- `DEVICE_OWNED_H2D` (正在上传)
|
|||
|
|
- `DEVICE_OWNED_COMPUTE` (正在计算)
|
|||
|
|
- `DEVICE_OWNED_D2H` (正在回传)
|
|||
|
|
- `RELEASED` (回传完毕,归还 Pool)
|
|||
|
|
|
|||
|
|
```mermaid
|
|||
|
|
stateDiagram-v2
|
|||
|
|
%% 状态定义
|
|||
|
|
state "HOST_OWNED<br/>(主机所有)" as HOST
|
|||
|
|
state "DEVICE_OWNED_H2D<br/>(传输中: H->D)" as H2D
|
|||
|
|
state "DEVICE_OWNED_COMPUTE<br/>(计算中: Kernel)" as COMPUTE
|
|||
|
|
state "DEVICE_OWNED_D2H<br/>(传输中: D->H)" as D2H
|
|||
|
|
state "RELEASED<br/>(待归还)" as RELEASED
|
|||
|
|
|
|||
|
|
%% 流程流转
|
|||
|
|
[*] --> HOST : 从 MemoryPool 申请
|
|||
|
|
|
|||
|
|
HOST --> H2D : I/O线程填充数据\n并调用 cudaMemcpyAsync
|
|||
|
|
note right of HOST
|
|||
|
|
此时数据位于页锁定内存
|
|||
|
|
CPU 写入完成
|
|||
|
|
end note
|
|||
|
|
|
|||
|
|
H2D --> COMPUTE : 记录 H2D_Event\nStreamWaitEvent
|
|||
|
|
note right of H2D
|
|||
|
|
DMA 引擎正在搬运
|
|||
|
|
CPU 不阻塞
|
|||
|
|
end note
|
|||
|
|
|
|||
|
|
COMPUTE --> D2H : Kernel 执行完毕\n自动触发 D2H
|
|||
|
|
note right of COMPUTE
|
|||
|
|
GPU 核心正在计算
|
|||
|
|
数据驻留显存
|
|||
|
|
end note
|
|||
|
|
|
|||
|
|
D2H --> RELEASED : D2H 完成回调\n或 Event 同步
|
|||
|
|
note right of D2H
|
|||
|
|
结果已写回 Host
|
|||
|
|
end note
|
|||
|
|
|
|||
|
|
RELEASED --> HOST : DataPacket 析构\n自动归还 Pool
|
|||
|
|
|
|||
|
|
RELEASED --> [*]
|
|||
|
|
```
|
|||
|
|
|
|||
|
|
-----
|
|||
|
|
|
|||
|
|
**下一步行动**:
|
|||
|
|
|
|||
|
|
我们已经定义了 **“怎么传”**(Pinned Memory)和 **“怎么调度”**(Async Streams)。
|
|||
|
|
接下来,我们需要解决 **“传给谁** 的问题,即 **2.2.3 NUMA 感知的内存亲和性控制**。考虑到飞腾 S5000C 的双路架构,如果数据传错了 CPU 节点,上述所有优化都会因为 QPI 总线瓶颈而大打折扣。
|
|||
|
|
|
|||
|
|
**提问**:您是否同意将 **“双流乒乓 (Double Stream Ping-Pong)”** 作为异步流水线的基线?确认后我们进入 2.2.3 NUMA 亲和性的讨论。
|