124 lines
7.2 KiB
Markdown
124 lines
7.2 KiB
Markdown
---
|
||
tags: []
|
||
date created: 星期四, 十一月 20日 2025, 11:20:35 晚上
|
||
date modified: 星期四, 十一月 20日 2025, 11:21:14 晚上
|
||
---
|
||
|
||
# 2.2.6 显存布局与对齐约束 (VRAM Layout & Alignment Constraints)
|
||
|
||
- **覆盖范围**：定义雷达数据立方体（Radar Data Cube）在显存中的物理排列格式。重点解决 **SoA (结构数组)** vs **AoS (数组结构)** 的选择、**Padding (填充)** 策略，以及适配 `cuFFT` / `cuBLAS` 库要求的复数存储格式。
|
||
|
||
## 一、 约束输入与对齐 (Constraints & Alignment)
|
||
|
||
基于 `02_信号处理模块设计.md` 和 GPU 硬件特性，我们要对齐以下约束：
|
||
|
||
1. **硬件合并访问 (Coalesced Access)**：GPU 读取显存的最佳模式是“一刀切”。同一个 Warp（32 个线程）必须访问一段**连续且对齐**的内存（通常是 128 字节）。如果数据是跳跃的（Strided），有效带宽会下降 80% 以上。
|
||
2. **雷达数据立方体特性**：数据具有三个维度：`[通道 (Channel)]`、`[脉冲 (Pulse)]`、`[距离门 (Range)]`。
|
||
3. **算法库约束**：
|
||
- **CoreX Math Libs (cuFFT)**：智铠重构版 `cuFFT` 通常要求输入数据为 **Interleaved Complex** (`float2` 或 `cuComplex`，即 `real, imag` 相邻) 或 **Split Complex** (`real[]`, `imag[]` 分离)。标准 CUDA 库倾向于 **Interleaved**。
|
||
4. **并行维度**：
|
||
- **脉冲压缩**：在 **距离门** 维度并行。这意味着“距离”维必须是内存中最连续的维度（Stride = 1）。
|
||
|
||
-----
|
||
|
||
## 二、 权衡分析与选项呈现 (Trade-off Matrix)
|
||
|
||
### 议题 1：复数数据格式 (Complex Number Format)
|
||
|
||
| 选项 | A. 交织存储 (Interleaved / AoS) **(推荐)** | B. 分离存储 (Split / SoA) |
|
||
| :--- | :--- | :--- |
|
||
| **格式** | `R I R I R I …` (`struct {float r, i}`) | `R R R …` / `I I I …` |
|
||
| **cuFFT 兼容性** | **原生支持**。`cufftExecC2C` 默认接受此格式。 | 需要使用 `cufftExecZ2Z` 并配置 stride，或者手动转换，稍显麻烦。 |
|
||
| **访存效率** | **高**。读取一个复数只需一次 64-bit 加载指令（`LD.E`）。 | **中**。读取一个复数需要两次 32-bit 加载指令，且地址相隔很远，增加指令发射压力。 |
|
||
| **结论** | **基线标准**。 | 不推荐，除非特定算法有强需求。 |
|
||
|
||
### 议题 2：数据立方体排列 (Data Cube Layout)
|
||
|
||
假设我们处理一个 `C` 通道、`P` 脉冲、`R` 距离门的数据块。
|
||
|
||
| 选项 | A. `[Channel][Pulse][Range]` (推荐) | B. `[Range][Pulse][Channel]` |
|
||
| :--- | :--- | :--- |
|
||
| **最内层维度** | **Range (距离)**。内存中连续存放 `R0, R1, R2…`。 | **Channel (通道)**。内存中连续存放 `C0, C1, C2…`。 |
|
||
| **脉冲压缩友好度** | **完美**。FFT 是针对 Range 做的，数据连续，读取效率 100%。 | **灾难**。FFT 需要读 Range 维，这里 Range 维跨度极大，导致严重的 TLB Miss 和非合并访问。 |
|
||
| **波束合成友好度** | **差**。DBF 需要跨通道计算。但在脉压之后做一次**转置**即可解决。 | **好**。 |
|
||
| **结论** | **基线标准**。符合“先脉压，后多普勒/DBF”的处理流。 | 仅适用于纯 DBF 前置的特殊雷达。 |
|
||
|
||
### 议题 3：行对齐与 Pitch (Padding Strategy)
|
||
|
||
显存是按“行”管理的。如果一行的字节数不是 256 字节的倍数，换行访问时就会错位，破坏对齐。
|
||
|
||
| 选项 | A. 紧凑排列 (Packed) | B. 对齐填充 (Pitched / Padded) **(推荐)** |
|
||
| :--- | :--- | :--- |
|
||
| **机制** | 数据紧挨着放。`Row1_End` 紧接 `Row2_Start`。 | 在每行末尾填充垃圾数据，使得 `Row_Stride` 是 256B 的倍数。 |
|
||
| **空间利用** | 100%。 | 略有浪费（\< 1%）。 |
|
||
| **访问性能** | **不稳定**。如果 `R` 不是 64 的倍数，第二行的起始地址就未对齐，导致 Warp 访问分裂，性能下降。 | **极致稳定**。确保每一行的起始地址都是对齐的，所有 Kernel 都能全速运行。 |
|
||
|
||
-----
|
||
|
||
## 三、 基线确立与实施规范
|
||
|
||
为了让 GPU 的吞吐量 KPI 达标，我们确立以下显存基线：
|
||
|
||
### 1\. 数据结构基线：Interleaved Complex (`float2`)
|
||
|
||
- **决策**：所有 I/Q 信号在显存中统一存储为 **`float2`** 类型（对应 `std::complex<float>`）。
|
||
- **规范**：
|
||
|
||
```cpp
|
||
// 严禁使用自定义结构体，直接用 CUDA 内置类型以确保编译器优化
|
||
using Complex = float2;
|
||
// 实部: .x, 虚部: .y
|
||
```
|
||
|
||
### 2\. 物理布局基线：`[Channel][Pulse][Range]` + Padding
|
||
|
||
- **决策**：采用 **Channel-Major, Pulse-Major, Range-Contiguous** 的布局，并对 **Range** 维度实施 **Pitch 对齐**。
|
||
- **二维视图**：将三维立方体视为一个巨大的二维矩阵：
|
||
- **宽度 (Width)**：`Range * sizeof(Complex)`
|
||
- **高度 (Height)**：`Channel * Pulse`
|
||
- **分配规范**：**必须**使用 `cudaMallocPitch` 分配显存，而不是 `cudaMalloc`。
|
||
- *论证*：`cudaMallocPitch` 会自动计算最佳的 `pitch`（字节步长），确保每一行（即每一个脉冲）的起始地址都满足硬件（Iluvatar MR-V100）的 256 字节对齐要求。
|
||
|
||
### 3\. 传输与访问规范
|
||
|
||
- **H2D 传输**：使用 `cudaMemcpy2DAsync`。
|
||
- 虽然数据源（Host Pinned Memory）可能是紧凑的，但目标（Device Memory）是有 Padding 的。`cudaMemcpy2D` 能自动处理这种“紧凑到对齐”的搬运。
|
||
- **Kernel 索引**：
|
||
|
||
```cpp
|
||
// Kernel 代码范式
|
||
// pitch 是字节单位，需要转为 float2 单位
|
||
int stride = pitch / sizeof(float2);
|
||
int p = blockIdx.y; // 脉冲索引 (或 C*P 的扁平索引)
|
||
int r = threadIdx.x + blockIdx.x * blockDim.x; // 距离门索引
|
||
|
||
if (r < Range) {
|
||
// 使用 stride 而不是 Range 来换行，跳过 Padding 区域
|
||
Complex data = input_data[p * stride + r];
|
||
// … 计算 …
|
||
}
|
||
```
|
||
|
||
## 总结：2.2 章节最终基线图谱
|
||
|
||
至此，我们完成了 **2.2 异构 DMA 与内存传输机制** 的全链路设计。从 Host 内存分配到 PCIe 传输，再到显存落地，每一环都扣上了“高性能”的锁扣。
|
||
|
||
| 环节 | 基线决策 | 关键配置/API |
|
||
| :--- | :--- | :--- |
|
||
| **2.2.1 内存分配** | **锁页内存 + UVA 映射** | `cudaMallocHost` + `cudaHostAllocMapped` |
|
||
| **2.2.2 传输调度** | **双流乒乓 (Double Buffering)** | 2x `cudaStream_t`, `cudaMemcpyAsync` |
|
||
| **2.2.3 物理位置** | **NUMA Node 1 本地化** | `numactl --cpunodebind=1 --membind=1` |
|
||
| **2.2.4 零拷贝** | **控制流 Zero-Copy，数据流 Explicit** | `WC` 标志 (控制流) vs `Async Copy` (数据流) |
|
||
| **2.2.5 传输粒度** | **≥ 64KB 块，256B 对齐** | 空间/时间双触发策略 |
|
||
| **2.2.6 显存布局** | **`[C][P][R]` 布局 + Pitch 对齐** | `cudaMallocPitch`, `float2` (Interleaved) |
|
||
|
||
-----
|
||
|
||
**下一步行动**：
|
||
|
||
2.2 章节已圆满完成。我们可以进入 **2.3 内部控制平面通信接口 (Internal Control Plane Interface - IPC)**。
|
||
|
||
虽然数据流（Data Plane）的问题解决了，但 Host 端的**任务调度器**如何优雅地指挥**信号处理模块**？如何不通过锁竞争就能实现高频的控制指令下发？这将是 2.3 讨论的重点。
|
||
|
||
**提问**：您是否确认 **2.2.6 的 Pitch 对齐与布局基线**？确认后我们结束 2.2 章节，进入 2.3。
|