> **适合人群**：有 C/C++ 基础，想搞加速卡编程的开发者 > **阅读时间**：约 20 分钟 > > **你将收获**：理解 DCU 硬件架构、内存体系、DMA 机制——为下篇编程实战打基础 > **📖 ****本文是系列上篇**，下篇已发布：[《DCU入门（下）：编程实战与性能优化》](./DCU%E5%85%A5%E9%97%A8%EF%BC%88%E4%B8%8B%EF%BC%89%EF%BC%9A%E7%BC%96%E7%A8%8B%E5%AE%9E%E6%88%98%E4%B8%8E%E6%80%A7%E8%83%BD%E4%BC%98%E5%8C%96.md) --- ## 目录 * [一、先聊点背景——DCU 是啥？](#%E4%B8%80%E5%85%88%E8%81%8A%E7%82%B9%E8%83%8C%E6%99%AFdcu-%E6%98%AF%E5%95%A5) * [二、DCU 的硬件——你的代码到底跑在什么东西上](#%E4%BA%8Cdcu-%E7%9A%84%E7%A1%AC%E4%BB%B6%E4%BD%A0%E7%9A%84%E4%BB%A3%E7%A0%81%E5%88%B0%E5%BA%95%E8%B7%91%E5%9C%A8%E4%BB%80%E4%B9%88%E4%B8%9C%E8%A5%BF%E4%B8%8A) * [2.1 整体架构：五大核心部件](#21-%E6%95%B4%E4%BD%93%E6%9E%B6%E6%9E%84%E4%BA%94%E5%A4%A7%E6%A0%B8%E5%BF%83%E9%83%A8%E4%BB%B6) * [2.2 CU 内部解剖：算力的最小单位](#22-cu-%E5%86%85%E9%83%A8%E8%A7%A3%E5%89%96%E7%AE%97%E5%8A%9B%E7%9A%84%E6%9C%80%E5%B0%8F%E5%8D%95%E4%BD%8D) * [2.3 Wavefront：理解 DCU 怎么干活的钥匙](#23-wavefront%E7%90%86%E8%A7%A3-dcu-%E6%80%8E%E4%B9%88%E5%B9%B2%E6%B4%BB%E7%9A%84%E9%92%A5%E5%8C%99) * [2.4 分支分歧——你的 if/else 为什么会让代码变慢](#24-%E5%88%86%E6%94%AF%E5%88%86%E6%AD%A7%E4%BD%A0%E7%9A%84-ifelse-%E4%B8%BA%E4%BB%80%E4%B9%88%E4%BC%9A%E8%AE%A9%E4%BB%A3%E7%A0%81%E5%8F%98%E6%85%A2) * [三、DCU 的内存——数据放哪，决定性能](#%E4%B8%89dcu-%E7%9A%84%E5%86%85%E5%AD%98%E6%95%B0%E6%8D%AE%E6%94%BE%E5%93%AA%E5%86%B3%E5%AE%9A%E6%80%A7%E8%83%BD) * [3.1 五级内存层次](#31-%E4%BA%94%E7%BA%A7%E5%86%85%E5%AD%98%E5%B1%82%E6%AC%A1) * [3.2 亲手测一遍——各内存性能对比实验](#32-%E4%BA%B2%E6%89%8B%E6%B5%8B%E4%B8%80%E9%81%8D%E5%90%84%E5%86%85%E5%AD%98%E6%80%A7%E8%83%BD%E5%AF%B9%E6%AF%94%E5%AE%9E%E9%AA%8C) * [3.3 Bank 冲突案例与解决方案](#33-bank-%E5%86%B2%E7%AA%81%E6%A1%88%E4%BE%8B%E4%B8%8E%E8%A7%A3%E5%86%B3%E6%96%B9%E6%A1%88) * [3.4 合并访问的本质](#34-%E5%90%88%E5%B9%B6%E8%AE%BF%E9%97%AE%E7%9A%84%E6%9C%AC%E8%B4%A8) * [四、HBM2——DCU 的"宽马路"显存](#%E5%9B%9Bhbm2dcu-%E7%9A%84%E5%AE%BD%E9%A9%AC%E8%B7%AF%E6%98%BE%E5%AD%98) * [五、DMA——CPU 与 DCU 之间的"自动搬运工"](#%E4%BA%94dmacpu-%E4%B8%8E-dcu-%E4%B9%8B%E9%97%B4%E7%9A%84%E8%87%AA%E5%8A%A8%E6%90%AC%E8%BF%90%E5%B7%A5) * [5.1 DMA 工作原理](#51-dma-%E5%B7%A5%E4%BD%9C%E5%8E%9F%E7%90%86) * [5.2 计算与传输重叠——Stream 实战](#52-%E8%AE%A1%E7%AE%97%E4%B8%8E%E4%BC%A0%E8%BE%93%E9%87%8D%E5%8F%A0stream-%E5%AE%9E%E6%88%98) --- ## 一、先聊点背景——DCU 是啥？ **如果你用过 NVIDIA 的 GPU，理解 DCU 就简单了——** **DCU = Deep Computing Unit（深度计算单元）**，由 **海光信息（Hygon）** 研发的国产 GPGPU（通用图形/计算处理器）。你可以把它理解成「国产版本的 GPU 加速卡」，专门干那些 CPU 搞不定的重型并行计算。 ### DCU 能干什么？ | **场景** | **具体应用** | | ---------------------- | ------------------------------------------------ | | **AI 训练/推理** | **千亿参数大模型训练、图像分类、目标检测** | | **科学计算** | **气象模拟、石油勘探、分子动力学、CFD** | | **大数据处理** | **基因测序、金融风险模型、推荐系统** | | **渲染与图形** | **3D 渲染、视频编解码** | ### 一个比喻帮你建立直觉 > **CPU 是一个"博士教授"**——会做微积分、写文章、跟人辩论，但一天只能干一件事。 > > **DCU 是一个"千人工厂流水线"**——每个工人只会做一件事（比如"加两个数"），但一千个人同时开工，一天能处理几百万个订单。 **这个差异决定了：****CPU 适合逻辑复杂的任务，DCU 适合数据量巨大的计算密集型任务。** --- ## 二、DCU 的硬件——你的代码到底跑在什么东西上 **不懂硬件就去写 DCU 代码，就像不了解发动机就去修车——能跑，但一调优就抓瞎。** ### 2.1 整体架构：五大核心部件 **一张 DCU 卡通过 ****PCI-E 接口** 插在服务器主板上，跟 CPU 配合干活： ** ** **记住这张表就行：** | **部件** | **一句话理解** | **关键数据** | | ------------------------------- | -------------------------------------------- | --------------------------------- | | **CU（计算单元）** | **DCU 的"肌肉"，负责一切数学运算** | **一块卡有几十到上百个 CU** | | **SIMD** | **用一条指令同时算多份数据** | **每 CU 含 4 个 SIMD** | | **HBM2 显存（全局内存）** | **DCU 的"工作台"，放要处理的数据** | **容量 16GB+，带宽 1TB/s** | | **DMA 引擎** | **硬件搬运工，CPU 下令它就自动搬数据** | **2 个引擎，互不干扰** | | **Cache（L1/L2）** | **缓存常用数据，减少慢速访存** | **硬件自动管理** | **关键认知**：DCU 有自己的显存，有独立的计算单元，有自己的 DMA——它是一个**高度独立的协处理器**，CPU 只负责给它"派活"和"收成果"。 ### 2.2 CU 内部解剖：算力的最小单位 **如果把 DCU 卡比作一栋写字楼，那 ****CU 就是其中一个办公室**。编程时，你的线程块（Block）被分配到一个 CU 上执行。我们聚焦一个 CU： ** ** **核心数据**： | **层级** | **数量** | **说明** | | --------------------------- | -------------------------- | ------------------------------------ | | **SIMD / CU** | **4 个** | **每个 SIMD 是独立指令流水线** | | **Wavefront / SIMD** | **10 个** | **硬件最多同时"记住"10 个 WF** | | **线程 / Wavefront** | **64 个** | **锁步执行（同一指令）** | | **矢量 ALU / SIMD** | **16 个**（DCU2 代） | **真正的计算单元** | | **最大活跃线程 / CU** | **2560** | **= 4 × 10 × 64** | ### 2.3 Wavefront：理解 DCU 怎么干活的钥匙 **在 NVIDIA 的世界里，调度单位叫 ****Warp（32 线程）**。在 DCU/AMD 的世界里，它叫 **Wavefront（64 线程）**——两者本质一样：**锁步执行的线程组**。 #### 什么是锁步执行（Lockstep）？ ``` DCU 执行指令: ADD R1, R2, R3 这意味着 64 个线程同时做加法： 线程 0: R1[0] = R2[0] + R3[0] 线程 1: R1[1] = R2[1] + R3[1] 线程 2: R1[2] = R2[2] + R3[2] ... 线程63: R1[63] = R2[63] + R3[63] 一条指令 → 64 路并行计算 → 能效极高 ``` #### 为什么需要 64 个线程一组？ **这跟硬件设计有关：****一个取指/译码单元就能驱动 64 个 ALU**。相比给每个 ALU 单独配一套控制逻辑，锁步设计更省电、芯片面积更小、频率可以更高。代价就是：遇到分支时会出问题。 ### 2.4 分支分歧——你的 if/else 为什么会让代码变慢 **这是 DCU 编程最容易踩的坑——****同一个 Wavefront 里的线程走了不同分支**。 **看个简单的例子：** ``` __global__ void divergent_kernel(float* data, int N) {    int idx = blockIdx.x * blockDim.x + threadIdx.x;        if (idx % 2 == 0) {        // 路径 A：偶数线程执行慢速计算        for (int i = 0; i < 10000; i++) {            data[idx] = sqrt(data[idx]) * 2.718f;       }   } else {        // 路径 B：奇数线程执行快速操作        data[idx] = data[idx] * 3.14f;   } } ``` **发生什么？** Wavefront 内 64 个线程的硬件执行过程： ``` 时间 → 步骤1: 先执行所有偶数线程(32个)的路径A ← 奇数线程被"屏蔽"（mask off），干等       [████████ 路径A: 10000次循环 █████████████████]       步骤2: 再执行所有奇数线程(32个)的路径B ← 偶数线程被屏蔽，干等       [█ 路径B: 1次乘法 █] 总时间 ≈ 路径A时间 + 路径B时间       ≈ 10000个循环的时间（路径B只占一小部分） ``` **理想情况**（所有线程走同一路径）： ``` 如果条件是 "threadIdx.x < 32"，而 Wavefront 内线程 0~63 是同组： 线程 0~31: 路径A ← 活跃 线程32~63: 路径B ← 被屏蔽   但这两个组各是半个 WF，依然会有分支分歧（如果是 Warp=32 就不会）！ ``` **性能影响**： | **分歧程度** | **执行时间** | **举例** | | -------------------------------- | ------------------ | ---------------------------------------------- | | **无分歧（全部相同路径）** | **1×** | `if (global_flag)`所有线程 global\_flag 相同 | | **2 路分歧** | **\~2×** | `if (threadIdx.x % 2)`偶数/奇数交替 | | **N 路分歧** | **\~N×** | `switch(threadIdx.x % 8)`8 种情况 | #### 怎么避免？ **策略一：让分支条件基于 blockIdx 而非 threadIdx** ``` // ✅ 好：不同 block 在不同 CU 上，不会互相阻塞 if (blockIdx.x % 2 == 0) {    // CU 0,2,4,... 的 block 走这里 } else {    // CU 1,3,5,... 的 block 走这里 } // 每个 CU 内部所有线程走同一路径 → 零分歧 ``` **策略二：用数学运算代替分支** ``` // ❌ 有分支 if (x > 0) y = x; else y = 0; // ✅ 无分支（三元运算符编译器可能帮优化，但不保证） y = fmaxf(x, 0.0f); // ✅✅ 无分支（确定无分支） float mask = (float)(x > 0);   // x>0 → 1.0, 否则 0.0 y = x * mask; ``` --- ## 三、DCU 的内存——数据放哪，决定性能 **搞 DCU 编程最核心的课题就是****内存管理**。数据放得不对，同一个算法性能可以差 10 倍。 ### 3.1 五级内存层次 ``` 速度（越来越快）             容量（越来越大）                               寄存器     ~1 周期     每线程最多 256 个     ← 最快的"草稿纸"   ↓ 共享内存   ~30 周期     每 CU 64KB           ← 线程块内的"公告板"   ↓ L1 Cache   ~30 周期     每 CU 几十 KB         ← 硬件自动缓存   ↓ L2 Cache   ~200 周期   全局几 MB             ← 跨 CU 共享   ↓ HBM2 显存   ~400 周期   16GB+               ← 数据的大仓库 ``` | **内存类型** | **你管还是硬件管** | **可见范围** | **适合放什么** | | ---------------------- | --------------------------------------------- | ---------------------- | -------------------------------- | | **寄存器** | **编译器管（你用局部变量暗示）** | **一个线程** | **循环变量、中间结果** | | **共享内存/LDS** | **你管**（`__shared__`关键字） | **一个线程块** | **线程间数据共享、缓冲区** | | **L1/L2 Cache** | **硬件管（你管不了）** | **L1=一个CU，L2=全局** | **有局部性的热点数据** | | **全局内存** | **你管**（`hipMalloc`/`hipMemcpy`） | **全局** | **输入/输出数组** | ### 3.2 亲手测一遍——各内存性能对比实验 **理论说完了，写个代码来****实测**这些内存的差异。这样你对"寄存器比全局内存快 400 倍"有真实感受。 ``` #include #include #include // ── 实验1：测试全局内存带宽 ── __global__ void global_memory_test(const float* __restrict__ in,                                     float* __restrict__ out, int N) {    int idx = blockIdx.x * blockDim.x + threadIdx.x;    if (idx < N) {        // 每次迭代：从全局内存读 → 做计算 → 写回全局内存        float val = in[idx];          // ~400 周期        val = val * 2.0f + 1.0f;        out[idx] = val;               // ~400 周期   } } // ── 实验2：利用寄存器减少访存 ── __global__ void register_optimized(const float* __restrict__ in,                                     float* __restrict__ out, int N) {    int idx = blockIdx.x * blockDim.x + threadIdx.x;    if (idx < N) {        float val = in[idx];          // 1 次全局读        // 以下操作全在寄存器内完成        float a = val * 2.0f;        float b = a + 1.0f;        float c = b * 0.5f;        float d = c + 3.0f;        float e = d * d;        out[idx] = e;                  // 1 次全局写   } } // ── 实验3：使用共享内存缓冲 ── #define BLOCK_SIZE 256 __global__ void shared_memory_test(const float* __restrict__ in,                                     float* __restrict__ out, int N) {    __shared__ float cache[BLOCK_SIZE];        int idx = blockIdx.x * blockDim.x + threadIdx.x;        // 线程块协作：每个线程从全局内存加载一个元素到共享内存    if (idx < N) cache[threadIdx.x] = in[idx];    else         cache[threadIdx.x] = 0.0f;        __syncthreads();  // 确保所有线程都加载完成        // 从共享内存（~30周期）读取，而非全局内存（~400周期）    float val = cache[threadIdx.x];    for (int i = 0; i < 1000; i++) {        val = val * 2.0f + 1.0f;  // 算在寄存器，读在共享内存   }    val += cache[(threadIdx.x + 128) % BLOCK_SIZE];  // 读邻居数据        if (idx < N) out[idx] = val; } int main() {    int N = 1 << 24;  // 16M 个元素    size_t size = N * sizeof(float);        float *h_in = new float[N];    float *h_out = new float[N];    for (int i = 0; i < N; i++) h_in[i] = i * 0.01f;        float *d_in, *d_out;    hipMalloc(&d_in, size);    hipMalloc(&d_out, size);        int block = BLOCK_SIZE;    int grid = (N + block - 1) / block;        hipMemcpy(d_in, h_in, size, hipMemcpyHostToDevice);        // 测实验1：全局内存版    auto t1 = std::chrono::high_resolution_clock::now();    global_memory_test<<>>(d_in, d_out, N);    hipDeviceSynchronize();    auto t2 = std::chrono::high_resolution_clock::now();    auto global_us = std::chrono::duration_cast(t2 - t1).count();        // 测实验2：寄存器优化版    t1 = std::chrono::high_resolution_clock::now();    register_optimized<<>>(d_in, d_out, N);    hipDeviceSynchronize();    t2 = std::chrono::high_resolution_clock::now();    auto reg_us = std::chrono::duration_cast(t2 - t1).count();        // 测实验3：共享内存版    t1 = std::chrono::high_resolution_clock::now();    shared_memory_test<<>>(d_in, d_out, N);    hipDeviceSynchronize();    t2 = std::chrono::high_resolution_clock::now();    auto shared_us = std::chrono::duration_cast(t2 - t1).count();        std::cout << "=== 内存性能对比 (16M 元素) ===" << std::endl;    std::cout << "全局内存频繁访存: " << global_us << " μs" << std::endl;    std::cout << "寄存器优化:       " << reg_us << " μs" << std::endl;    std::cout << "共享内存缓冲:     " << shared_us << " μs" << std::endl;        hipFree(d_in);    hipFree(d_out);    delete[] h_in;    delete[] h_out;    return 0; } ``` **预期结果**（实际数字取决于 DCU 型号）： ``` === 内存性能对比 (16M 元素) === 全局内存频繁访存: 450 μs     ← 每次都从 HBM2 读写，最慢 寄存器优化:       180 μs     ← 中间计算全在寄存器，比上面快 2.5× 共享内存缓冲:     120 μs     ← 共享内存降低延迟，比全局快 ~4× ``` ### 3.3 Bank 冲突案例与解决方案 **共享内存虽快，但有个重要的****性能陷阱**——Bank 冲突。 #### Bank 是什么？ ``` 共享内存分区示意图： Bank 0 Bank 1 Bank 2 ... Bank 31 [4字节] [4字节] [4字节]       [4字节] [4字节] [4字节] [4字节]       [4字节]   ...                                 ← 共 32 个 Bank ``` **32 个 Bank 能****同时**返回 32 个数据。但如果多个线程访问**同一个 Bank**，就会产生冲突，请求被串行化。 #### 冲突重现与修复 ``` // ── 示例：从共享内存读取数据 ── __global__ void bank_conflict_demo(float* input, float* output, int N) {    __shared__ float smem[32][32];        int idx = blockIdx.x * blockDim.x + threadIdx.x;    int tx = threadIdx.x % 32;    int ty = threadIdx.x / 32;        // 加载数据到共享内存    for (int i = 0; i < 32; i++) {        smem[ty][tx] = input[blockIdx.x * 1024 + ty * 32 + tx];   }    __syncthreads();    // ─── 场景1：无冲突 (stride=1) ───    // 线程0→Bank0, 线程1→Bank1, ... 线程31→Bank31    // 所有 Bank 同时服务 → 完美    float val1 = smem[0][tx];  // 或 smem[threadIdx.x];    // ─── 场景2：32路冲突 (stride=32) ───    // 所有线程访问 smem[0][0], smem[1][0], smem[2][0], ...    // smem[0][0] 在 Bank0, smem[1][0] 也在 Bank0(因为每行32个元素)    // → 32个线程全部竞争 Bank0 → 32路冲突 → 性能暴跌32倍!    float val2 = smem[tx][0];    // ─── 场景3：加 padding 解决冲突 ───    __shared__ float smem_padded[32][33];  // 多一列!    // 现在 smem_padded[0][0] → Bank0    //     smem_padded[1][0] → Bank1 (因为每行 33 个元素)    //   → 无冲突! } // ─── 完整示例：矩阵转置中的 Bank 冲突 ─── #define TILE 32 #define TILE_PAD 33  // 加 padding // V1：有 Bank 冲突的版本（列读取） __global__ void transpose_naive(float* __restrict__ out,                                 const float* __restrict__ in,                                 int rows, int cols) {    __shared__ float tile[TILE][TILE];  // 32×32 = 1024 个 float        int x = blockIdx.x * TILE + threadIdx.x;  // 列    int y = blockIdx.y * TILE + threadIdx.y;  // 行        // 从全局内存加载到共享内存（行读取 → 无冲突）    tile[threadIdx.y][threadIdx.x] = in[y * cols + x];    __syncthreads();        // 从共享内存读取后写回全局（列读取 → Bank 冲突！）    // 线程(0,0) 读 tile[0][0] → Bank0    // 线程(1,0) 读 tile[1][0] → Bank0 ← 冲突！    int out_x = blockIdx.y * TILE + threadIdx.x;    int out_y = blockIdx.x * TILE + threadIdx.y;    out[out_y * rows + out_x] = tile[threadIdx.x][threadIdx.y]; } // V2：消除 Bank 冲突的版本（padding） __global__ void transpose_optimized(float* __restrict__ out,                                     const float* __restrict__ in,                                     int rows, int cols) {    __shared__ float tile[TILE][TILE_PAD];  // 32×33，消除冲突        int x = blockIdx.x * TILE + threadIdx.x;    int y = blockIdx.y * TILE + threadIdx.y;        tile[threadIdx.y][threadIdx.x] = in[y * cols + x];    __syncthreads();        int out_x = blockIdx.y * TILE + threadIdx.x;    int out_y = blockIdx.x * TILE + threadIdx.y;    out[out_y * rows + out_x] = tile[threadIdx.x][threadIdx.y]; } ``` ### 3.4 合并访问的本质 **合并访问（Coalesced Access）** 是 DCU 编程最重要的优化原则之一。 ``` 原则：同一 Wavefront 的 64 个线程访问相邻的内存地址 ✅ 合并访问 (coalesced)：   线程0→地址A 线程1→A+4 线程2→A+8 ... 线程63→A+252   → 一次内存事务，传输 256 字节（64×4B），利用率 100% ❌ 非合并访问 (strided)：   线程0→A+0   线程1→A+256 线程2→A+512 ...     → 至少 64 次内存事务（如果是 stride=64），利用率 1/64 ❌ 随机访问：   线程0→随机地址 线程1→随机地址 ...   → 每次都 cache miss，几十次单独事务 ``` **验证代码**： ``` // 合并 vs 非合并访问对比 __global__ void coalesced_vs_strided(    const float* __restrict__ in, float* __restrict__ out,    int N, float* __restrict__ out2) {    int idx = blockIdx.x * blockDim.x + threadIdx.x;        // ✅ 合并访问    if (idx < N) out[idx] = in[idx] * 2.0f;        // ❌ 非合并访问：相邻线程读的元素间隔 64 个    //   线程0→in[0], 线程1→in[64], 线程2→in[128]...    int gid = blockIdx.x * blockDim.x + threadIdx.x * 64;    if (gid < N) out2[gid] = in[gid] * 2.0f; } ``` --- ## 四、HBM2——DCU 的"宽马路"显存 **DCU 使用 ****HBM2（High Bandwidth Memory 第二代）**，而非普通 DDR 内存。理解它能帮你理解为什么某些访问模式快、某些慢。 ### HBM2 怎么做到 1TB/s 的？ | **对比** | **DDR4 (CPU 内存)** | **GDDR6 (普通显卡)** | **HBM2 (DCU 高端显卡)** | | -------------- | ------------------------- | -------------------------- | ----------------------------- | | **位宽** | **64 bit** | **32 bit/颗粒** | **1024 bit** | | **频率** | **\~3200 MHz** | **\~14000 MHz** | **\~2000 MHz** | | **带宽** | **\~25 GB/s** | **\~400 GB/s** | **\~1000 GB/s** | | **功耗** | **中** | **高** | **低**（频率低+电压低） | | **成本** | **低** | **中** | **高** | **HBM2 的策略：****不高频 → 加大位宽**。用 1024 条数据线同时传，比 DDR4 的 64 条宽了 16 倍，一个时钟周期就能传 128 字节。 ### 物理原理：3D 堆叠 ``` 传统内存（平面布局）：           HBM2（3D堆叠）： ┌──────────────────┐           ┌─────────┐ ← DRAM Die 4 │ 1颗 DRAM 芯片   │             ├─────────┤ ← DRAM Die 3 └──────────────────┘           ├─────────┤ ← DRAM Die 2                   ← PCB走线链接 ├─────────┤ ← DRAM Die 1 ┌──────────────────┐           ├─────────┤ ← 逻辑控制 Die │     GPU Die     │           │ TSV直连 │ (垂直硅通孔) └──────────────────┘           └────┬────┘                                   │ Silicon Interposer                               ┌────┴────┐                               │ DCU Die │ ← 紧邻，信号路径极短                               └─────────┘ ``` **垂直堆叠 + 硅中介层：数据传输距离从厘米级降到毫米级 → 延迟降低 + 功耗降低。** ### 对我们写代码意味着什么？ 1. **位宽巨大 → 连续访问收益极高**：一次事务就能传 128 字节，合并访问事半功倍 2. **延迟依然高（\~400 周期）→ 要大量线程掩藏延迟** 3. **容量有限（16GB）→ 大模型要用数据分块（tiling）** --- ## 五、DMA——CPU 与 DCU 之间的"自动搬运工" ### 5.1 DMA 工作原理 **没有 DMA 的时候，CPU 亲自搬数据：** ``` CPU: 从内存读1字节 → 写到PCI-E寄存器 → 再读1字节 → 再写...     ↑ 全程占用 CPU，别的活都得等 ``` **有了 DMA，CPU 只需要下一道命令：** ``` CPU: "DMA，帮我把这 4GB 数据从主存搬到 DCU 显存" DMA: 收到 → 自己搬 → 搬完了 → 发中断通知 CPU CPU: (在 DMA 搬运期间可以干别的事) ``` **DCU 有 ****2 个 DMA 引擎**： * **Engine 0：CPU ↔ DCU 数据传输** * **Engine 1：DCU ↔ DCU 数据传输（多卡场景）** ### 5.2 计算与传输重叠——Stream 实战 **最大的优化机会**：让 DMA 搬运数据和 DCU 计算**同时进行**。 ``` #include #include int main() {    const int N = 1 << 24;    const int chunk_size = N / 4;    const size_t chunk_bytes = chunk_size * sizeof(float);        // 创建两个 stream    hipStream_t stream1, stream2;    hipStreamCreate(&stream1);    hipStreamCreate(&stream2);        // 主机内存    float *h_a = new float[N];    float *h_b = new float[N];    float *h_c = new float[N];    for (int i = 0; i < N; i++) { h_a[i] = i; h_b[i] = 2*i; }        // DCU 显存    float *d_a, *d_b, *d_c;    hipMalloc(&d_a, N * sizeof(float));    hipMalloc(&d_b, N * sizeof(float));    hipMalloc(&d_c, N * sizeof(float));        int block = 256;    int grid = (chunk_size + block - 1) / block;        // ─── 串行版本（对比基准） ───    hipMemcpy(d_a, h_a, N * sizeof(float), hipMemcpyHostToDevice);    hipMemcpy(d_b, h_b, N * sizeof(float), hipMemcpyHostToDevice);    hipLaunchKernelGGL(vectorAdd, grid * 4, block, 0, 0, d_a, d_b, d_c, N);    hipMemcpy(h_c, d_c, N * sizeof(float), hipMemcpyDeviceToHost);    // DMA kernel1 DMA — 全部串行        // ─── 异步流水线版本（高性能） ───    // 分块 0: 先启动 DMA (stream1)    hipMemcpyAsync(d_a, h_a, chunk_bytes, hipMemcpyHostToDevice, stream1);    hipMemcpyAsync(d_b, h_b, chunk_bytes, hipMemcpyHostToDevice, stream1);        // 分块 1: 先启动 DMA (stream2)，与分块0的kernel同时进行    hipMemcpyAsync(d_a + chunk_size, h_a + chunk_size, chunk_bytes,                   hipMemcpyHostToDevice, stream2);    hipMemcpyAsync(d_b + chunk_size, h_b + chunk_size, chunk_bytes,                   hipMemcpyHostToDevice, stream2);        // 分块 0 的计算可以和分块 1 的 DMA 重叠    hipLaunchKernelGGL(vectorAdd, grid, block, 0, stream1,                       d_a, d_b, d_c, chunk_size);        // 分块 1 的计算和分块 2 的 DMA 重叠    hipLaunchKernelGGL(vectorAdd, grid, block, 0, stream2,                       d_a + chunk_size, d_b + chunk_size, d_c + chunk_size, chunk_size);        // 分块 2 的 DMA    hipMemcpyAsync(d_a + 2*chunk_size, h_a + 2*chunk_size, chunk_bytes,                   hipMemcpyHostToDevice, stream1);    hipMemcpyAsync(d_b + 2*chunk_size, h_b + 2*chunk_size, chunk_bytes,                   hipMemcpyHostToDevice, stream1);        // 分块 2 的计算    hipLaunchKernelGGL(vectorAdd, grid, block, 0, stream1,                       d_a + 2*chunk_size, d_b + 2*chunk_size, d_c + 2*chunk_size, chunk_size);        // ... 以此类推        hipDeviceSynchronize();  // 等所有 stream 完成        hipFree(d_a); hipFree(d_b); hipFree(d_c);    delete[] h_a; delete[] h_b; delete[] h_c;        hipStreamDestroy(stream1);    hipStreamDestroy(stream2);    return 0; } ``` **流水线效果**： ``` 时间 → 串行: |──── DMA大块 ────|──── 计算大块 ────|──── 回传大块 ────|       ↑ 计算等着 DMA，浪费 流水线: |─ DMA块0 ─|─ DMA块1 ─|─ DMA块2 ─|─ ...       |   Kernel块0   |   Kernel块1   |   ↑ 重叠，互相不等待 ``` --- > **📖 ****下篇已发布**：[《DCU入门（下）：编程实战与性能优化》](./DCU%E5%85%A5%E9%97%A8%EF%BC%88%E4%B8%8B%EF%BC%89%EF%BC%9A%E7%BC%96%E7%A8%8B%E5%AE%9E%E6%88%98%E4%B8%8E%E6%80%A7%E8%83%BD%E4%BC%98%E5%8C%96.md) — 从向量加法到矩阵乘法 V4，手把手写 HIP 代码，内含 50 倍性能优化全记录。 **如果觉得有帮助，欢迎 ****点赞 + 收藏**，有问题评论区见。