# DCU K100 节点初探：从硬件配置到性能调优全记录 > **摘要**：本文以海光 DCU K100 加速卡为目标平台，从开箱即用的硬件环境出发，系统梳理了 ROCm 软件栈、NUMA 拓扑与 PCIe 带宽分析、Hipprof 性能剖析工具以及 TensorCore 编程模型。文章结合实测数据，给出了带宽利用率分析和 Kernel 性能优化方向，旨在为首次接触 DCU 平台的开发者提供一份可操作的上手指南。 --- ## 1. 引言 随着国产异构计算平台的快速发展，海光（Hygon）推出的 DCU（Deep Computing Unit）系列加速卡在科学计算、分子动力学模拟等领域逐渐崭露头角。DCU 兼容 AMD ROCm（Radeon Open Compute platforM）开放式软件栈，开发者可以沿用 HIP（Heterogeneous-Compute Interface for Portability）编程模型，在较低的迁移成本下将现有 GPU 应用移植到国产平台。 本文将基于一台搭载 8 张 DCU K100 的服务器节点，逐层展开：从最基础的硬件信息采集，到 NUMA 拓扑与 PCIe 带宽的实测分析，再到 Hipprof 性能剖析工具的使用，最后探讨 TensorCore 的编程调用方式。希望这篇记录能帮助读者快速建立起对 DCU 平台的整体认知。 --- ## 2. 硬件环境概览 ### 2.1 CPU 与 NUMA 拓扑 登录节点后，首先使用 `lscpu` 命令查看 CPU 信息： ``` Architecture: x86_64 CPU(s): 64 Model name: Hygon C86 7381 32-core Processor Socket(s): 2 NUMA node(s): 8 ```  节点搭载两颗 **Hygon C86 7381** 处理器，每颗 32 核，总计 64 物理核心。值得注意的是，系统被划分为 **8 个 NUMA 节点**——这一拓扑结构将直接影响后续 DCU 加速卡的亲和性绑定和跨 NUMA 通信性能。 > **关键认知**：在 NUMA 架构下，CPU 访问本地内存与远端内存的延迟差异显著。同理，DCU 与 CPU 之间的数据传输也应尽量保持在同一个 NUMA 节点内，避免跨节点 PCIe 桥接带来的性能损耗。 ### 2.2 DCU 硬件信息 使用 `rocminfo` 命令可以列举 ROCm 平台下所有可见的硬件设备： ``` Agent 8 (DCU K100): Name: gfx926 Max Threads Per Block: 1024 ... ```  关键参数提炼如下： | 参数 | 值 | 影响 | | ----------------------- | --------------------- | ------------------------------------ | | **型号** | DCU K100 | — | | **架构代号** | gfx926 | 编译时需指定 `--offload-arch=gfx926` | | **单 Block 最大线程数** | 1024 | 决定 Kernel 的 work-group size 上限 | | **PCIe 接口** | PCIe 4.0 ×16, 16 GT/s | 决定主机与设备间的数据传输上限 | ### 2.3 ROCm 软件栈简介 ROCm 是 AMD 发起的开放式 GPU 计算平台，对标 NVIDIA CUDA，其核心组件包括： - **HIP**：异构计算可移植接口，提供与 CUDA 语法高度相似的 C++ API，支持 `hipify` 工具自动转换 CUDA 代码。 - **ROCk / ROCt**：内核驱动与运行时，负责设备管理、内存分配、Kernel 分发。 - **rocBLAS / rocFFT / rocSOLVER**：对标 cuBLAS / cuFFT / cuSOLVER 的基础数学库。 - **Hipprof**：DCU 平台的性能剖析工具（详见第 4 节）。 这些组件共同构成了一套完整的"编译器—运行时—数学库—调试工具"工具链。 --- ## 3. 带宽测试与 NUMA 拓扑分析 带宽是衡量异构计算平台数据传输能力的核心指标。本节通过 `rocm-bandwidth-test` 工具进行系统化的点对点带宽测试，并结合 NUMA 拓扑给出分析。 ### 3.1 测试环境与方法 测试节点中共有 **16 个设备**：编号 `0~7` 为 CPU 端（8 个 NUMA 节点的处理器侧），编号 `8~15` 为 8 张 DCU K100 加速卡。 使用如下命令进行测试： ```bash # 所有设备间点对点单向带宽 rocm_bandwidth_test -a # 所有设备间点对点双向带宽 rocm_bandwidth_test -A # 指定设备对之间的延迟测试 rocm_bandwidth_test -s -d -l ``` ### 3.2 单向带宽结果分析 首先，通过 DCU 之间的"NUMA 距离"图可以直观看出设备的物理分组关系：  从上图可以看出 8~11 号为一组，12~15 号为另一组。下表展示了 8 张 DCU（设备 8~15）之间的单向点对点带宽（单位：GB/s）： | 分组 | DCU 8~11（组内） | DCU 12~15（组内） | 跨组 | | ------------ | ---------------- | ----------------- | -------- | | **单向带宽** | ~28 GB/s | ~28 GB/s | ~11 GB/s |  这一结果恰好与 `rocm-smi --showtopo` 揭示的 NUMA 归属一致： ``` DCU 0~3 (设备 8~11) → 从属 NUMA 3 DCU 4~7 (设备 12~15) → 从属 NUMA 7 ``` **结论**：同一 NUMA 组内的 DCU 间通信带宽约为 28 GB/s，跨组带宽骤降至 11 GB/s，降幅达 **60%**。 ### 3.3 双向带宽与理论峰值对比 双向带宽测试结果： | 分组 | 组内 | 跨组 | | ------------ | -------- | -------- | | **双向带宽** | ~55 GB/s | ~22 GB/s |  进一步通过 `rocm-smi --showtopo` 查看拓扑结构：  可以看到 0~3 号 DCU 从属于 NUMA 3，4~7 号 DCU 从属于 NUMA 7，所有 DCU 均通过 PCIe 连接。 再通过 `rocm-smi --detail` 确认 PCIe 规格： ``` PCIe Version: 4.0 PCIe Speed: 16 GT/s PCIe Width: ×16 ``` 参考 PCIe 4.0 ×16 的理论带宽（编码方式为 128b/130b）：  | 指标 | 理论峰值 | 实测值 | 利用率 | | ------------ | --------- | ------- | --------- | | **单向带宽** | 31.5 GB/s | 28 GB/s | **88.9%** | | **双向带宽** | 63.0 GB/s | 55 GB/s | **87.0%** | > 实测带宽达到理论峰值的 **87%–89%**，属于 PCIe 4.0 ×16 的正常水平。剩余开销主要来自 PCIe 协议层（TLP 头、DLLP、流控等）以及 ROCm 运行时的软件开销。 ### 3.4 实践建议 基于上述分析，在应用部署时应注意： 1. **进程绑定**：将 MPI 进程绑定到与目标 DCU 相同 NUMA 节点的 CPU 核心上，避免跨 NUMA 数据传输。 2. **数据传输聚合**：尽量减少跨组 DCU 之间的直接数据交换，优先在组内完成归约等集合通信操作。 3. **PCIe 带宽利用**：通过流水线（Pipelining）技术将数据传输与计算重叠，进一步隐藏 PCIe 延迟。 --- ## 4. 性能剖析工具：Hipprof Hipprof 是 DCU 平台提供的性能分析工具，具备三大核心能力： - **可视化性能时间线**：展示各 Kernel 和 API 调用的时间分布。 - **HIP API 追踪**：记录所有 HIP API 调用及其参数和耗时。 - **PMC 硬件计数器采集**：读取 GPU 内部的性能监控计数器（Performance Monitor Counter）。 ### 4.1 硬件计数器采集 以下命令以 8 进程并行运行目标程序，并采集 PMC 数据： ```bash mpirun -np 8 hipprof --pmc ./clmd-mpi-dcu ``` 下图展示了描述符导数计算 Kernel 的 PMC 性能数据：  从图中可知该 Kernel：1. 寄存器使用较少（vgpr、sgpr）；2. ALU 利用率较低（7.8%）。对于上述两点，可以增加循环展开，同时优化访存以增加计算密度。 | 参数 | 当前值 | 诊断 | | ------------------------------ | ------ | ------------------------------------------ | | **vgpr count** | 较少 | 每个线程的矢量寄存器压力不大 | | **sgpr count** | 较少 | 标量寄存器使用同样较轻 | | **processed ALU instructions** | 7.8% | ALU 利用率低，处于访存瓶颈（Memory-Bound） | **优化方向**： - **增加循环展开**：减少循环控制开销，提高指令级并行度。 - **优化访存模式**：使用共享内存（LDS）缓存复用数据，提升计算密度（Compute Intensity）。 - **合并内存访问**：确保相邻线程访问相邻的内存地址（Coalesced Access），充分利用内存带宽。 ### 4.2 可视化性能分析 ```bash mpirun -np 8 hipprof --hip-trace --mpi-trace --group-stream ./clmd-mpi-dcu ``` 该命令会生成包含 HIP API 轨迹、MPI 通信轨迹、以及 Kernel 执行时间线的可视化报告：  可直观定位： - **Kernel 启动延迟**：API 调用到 Kernel 实际执行的时间间隔。 - **通信与计算重叠度**：MPI 通信是否与 GPU 计算有效重叠。 - **Kernel 串行化问题**：多个 Kernel 是否被意外地分配到同一 Stream 导致串行执行。 ### 4.3 Hipprof 参数速查表 #### Kernel 调度信息 | 参数 | 含义 | | ------------------------- | --------------------------------------------------- | | **kernel dispatch index** | Kernel 分发序号——程序启动后驱动分发的第 N 个 Kernel | | **gpu id** | 目标 GPU 编号——多卡环境下用于定位 Kernel 运行位置 | | **dispatch queue id** | 分发队列 ID——驱动内部的任务流硬件/软件标识符 | | **dispatch queue index** | 分发队列索引——Kernel 在当前队列中的排队位置 | | **thread id** | 主机端线程 ID——发出 Kernel 启动命令的 CPU 线程 | #### Kernel 资源占用 | 参数 | 含义 | | ---------------------- | ------------------------------------------------------------ | | **grid size** | 网格总线程数 = `Block 数 × work-group size` | | **work group size** | 每个 Block（线程块）内的线程数 | | **shared memory size** | 每个 Block 申请的 LDS 大小（字节） | | **scratch size** | 寄存器溢出（Spill）导致的暂存区大小——**非 0 时性能严重受损** | | **vgpr count** | **每线程**占用的矢量通用寄存器数 | | **sgpr count** | **每 Wavefront**占用的标量通用寄存器数 | | **fbarrier count** | 硬件栅栏槽位/状态计数 | | **signal handle** | 异步操作间同步与通信的内部句柄指针 | #### Kernel 性能指标 | 参数 | 含义 | | ------------------------------ | ------------------------------------------------ | | **kernel time** | GPU 上纯 Kernel 运行时间（秒） | | **processed ALU instructions** | ALU 指令执行时间占比——**极低通常意味着访存瓶颈** | | **performance** | 算力吞吐量（Gflops） | --- ## 5. DCU TensorCore 编程实战 DCU K100 集成了矩阵计算加速单元（TensorCore），可在一条指令内完成小规模矩阵乘加运算，大幅提升 AI 和 HPC 场景的吞吐量。本节从标准 GPU Kernel 出发，逐步演进到高性能 TensorCore 实现，完整呈现开发过程中的关键代码与优化思路。 ### 5.1 基础知识 #### 数据类型与宏定义 在编写 TensorCore 代码前，需要先定义向量化的数据类型和累加器清零宏： ```cpp // 4 元素向量类型：用于 TensorCore 累加器 typedef float float4_ __attribute__((ext_vector_type(4))); typedef double double4_ __attribute__((ext_vector_type(4))); typedef double double2_ __attribute__((ext_vector_type(2))); // 累加器清零宏 #define CLEAR_ACC(x4) \ x4.x = 0.0; \ x4.y = 0.0; \ x4.z = 0.0; \ x4.w = 0.0; ``` `float4_` 是 TensorCore 的核心数据结构——每次矩阵乘加指令的输出和累加都通过它完成，每个线程持有 `float4_` 中的 4 个元素。 #### TensorCore 计算模型 TensorCore 的计算以 **Wavefront（64 线程）** 为基本单位： - 每个 Warp 中的 **64 个线程**协同完成一次矩阵乘加：**D = A × B + C** - 使用 TensorCore 的 Kernel，其 Block 大小至少为 **64**（且应是 64 的倍数） - 单条指令完成 **16×16×4** 的矩阵乘（M=16, N=16, K=4），每个线程得到 4 个结果值，所有线程合计产出 **4 × 64 = 256** 个值 - 数据分布如图所示：T0表示wave中0号线程，以此类推T63表示wave中63号线程。T0中V0表示0号线程拿到的数据v0；T0中V1表示0号线程拿到的数据v1。64个数据的内存存储按照M/N的方向。如上图所示线程0～15号取M/N方向的16个数据；16～31号线程取第k+1列/行的16个数据；32～47号线程取第k+2列/行的16个数据；48～63号线程取第k+3列/行16个数据。64个数据分配到64个线程上，单个线程取一个数据。 - 输出数据大小为16x16的矩阵块。单个线程存储的数据大小为16x16/64个线程=4个数据。如图所示T0线程 输出的数据为T0-V0、T0-V1、T0-V2、T0-V3。4个输出数据不连续，每4个数据存放一个数据。  --- ### 5.2 起点：标准 HIP 矩阵乘 Kernel 在引入 TensorCore 之前，先看一个标准的矩阵乘 GPU Kernel，作为对比基准。矩阵 A 为 **列主序（col-major）**，矩阵 B 为 **行主序（row-major）**： ```cpp // A: col-major [M×K], B: row-major [K×N], C/D: col-major [M×N] __global__ void MatMulAdd_kernel2( data_type *A, data_type *B, data_type *C, data_type *D, int M, int N, int K) { int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; if (row < M && col < N) { data_type sum = 0; for (int k = 0; k < K; k++) { sum += A[row + k * M] * B[k * N + col]; } D[row + col * M] = sum + C[row + col * M]; } } // Host 端调用 void MatMulAdd_gpu(data_type *A, data_type *B, data_type *C, data_type *D, int M, int N, int K) { dim3 block(16, 16); dim3 grid((N + 15) / 16, (M + 15) / 16); hipLaunchKernelGGL(MatMulAdd_kernel2, grid, block, 0, 0, A, B, C, D, M, N, K); } ``` 这个 Kernel 每个线程负责计算输出矩阵的一个元素，通过内层循环遍历 K 维度完成点积。其计算密度受限于全局内存带宽，是典型的 **Memory-Bound** 实现。 --- ### 5.3 两种调用方式 以 `M=16, N=16, K=4` 的矩阵乘为例，DCU TensorCore 提供两种调用方式。 #### 方式一：Builtin 内建函数 Builtin 函数接收三个参数：映射矩阵 A 元素的变量、映射矩阵 B 元素的变量、储存结果的 `float4_` 变量：   ```cpp float4_ D_d; CLEAR_ACC(D_d); // __builtin_amdgcn_mmac_f32_16x16x4f32: 单精度矩阵乘加 // M=16, N=16, K=4, 结果累加到 D_d D_d = __builtin_amdgcn_mmac_f32_16x16x4f32(A_val, B_val, D_d); ``` **Tips**：builtin 的正确指令名是 `__builtin_amdgcn_mmac_f32_16x16x4f32`（注意 `mmac` 而非 `mfma`，前者是 AMD CDNA 架构的矩阵乘加指令）。 #### 方式二：内联汇编（Inline Assembly） 当需要更精细的寄存器控制或使用双精度时，采用内联汇编：   ```cpp float4_ D_d; CLEAR_ACC(D_d); // 单精度: v_mmac_16x16x4_f32 asm volatile( "s_waitcnt lgkmcnt(0)\n\t" "v_mmac_16x16x4_f32 %0, %1, %2, %0\n\t" : "+v"(D_d) // %0: 累加器（输入+输出） : "v"(A_val), "v"(B_val) // %1,%2: 矩阵 A、B 元素 ); // 双精度: v_mmac_16x16x4_f64 asm volatile( "s_waitcnt lgkmcnt(0)\n\t" "v_mmac_16x16x4_f64 %0, %1, %2, %0\n\t" : "+v"(D_d) : "v"(A_val), "v"(B_val) ); ``` > **关于 `s_waitcnt lgkmcnt(0)`**： > > `s_waitcnt` 全称 **S**calar **Wait** **C**ou**nt**（标量等待计数），是一条标量单元控制指令，作用是让 GPU 执行单元"停一下"，直到指定的异步操作全部完成。 > > `lgkmcnt` 是三个监控目标的缩写组合： > > - **L**（LDS）— 追踪从共享内存读取数据的指令进度 > - **G**（Global / Generic）— 追踪从全局内存读取数据的指令进度 > - **KM**（K-cache / Constant / Message）— 追踪从常量缓存或标量消息通路读取数据的进度 > > `lgkmcnt(0)` 表示等待计数器归零——即所有被追踪的访存操作均已完成，数据已真正落入寄存器，此时发射依赖这些数据的 `v_mmac_*` 计算指令才是安全的。**省略这一条指令可能导致计算单元读到"半生不熟"的旧数据，引发数值错误且极难排查。** > > **选择建议**：优先使用 Builtin 方式，代码简洁可移植；双精度场景或需要精确排布指令时使用内联汇编。 --- ### 5.4 最小 TensorCore Kernel：16×16×4 下面是最精简的 TensorCore Kernel——每个 Block 只有 64 个线程，全局内存只加载一次 A 和 B（K=4），然后直接调用 TensorCore 指令： ```cpp __global__ void Matrix4_dot_gpu_tc_kernel( data_type *A, data_type *B, data_type *C, data_type *D) { int ip = blockIdx.x * blockDim.x + threadIdx.x; // 0~63 __shared__ data_type A_s[64]; __shared__ data_type B_s[64]; // 加载到共享内存（coalesced access） A_s[ip] = A[ip]; B_s[ip] = B[ip]; __syncthreads(); float4_ D_d; CLEAR_ACC(D_d); #ifdef USE_FLOAT D_d = __builtin_amdgcn_mmac_f32_16x16x4f32(A_s[ip], B_s[ip], D_d); #endif int row = ip % 16; int col = ip / 16; // 将 float4_ 的 4 个元素写回全局内存 D[row * 16 + col + 0 * 4] = D_d.x; D[row * 16 + col + 1 * 4] = D_d.y; D[row * 16 + col + 2 * 4] = D_d.z; D[row * 16 + col + 3 * 4] = D_d.w; } // Host 端调用：只需 1 个 Block，64 个线程 dim3 blocks(64, 1, 1); dim3 grids(1, 1, 1); hipLaunchKernelGGL(Matrix4_dot_gpu_tc_kernel, grids, blocks, 0, 0, A, B, C, D); ``` 这个 Kernel 清晰地展示了 TensorCore 的编程范式：**加载到共享内存 → 调用一条 TensorCore 指令 → 写回结果**。但它只处理了 K=4 的固定问题，无法发挥实际性能。 --- ### 5.5 K 维度展开：从 16×16×4 到 16×16×32 实际场景中 K 通常远大于 4。通过循环遍历 K 维度，每次处理 K=4 的一个分片（tile），并将结果累加到同一组 `float4_` 中： ```cpp __global__ void vmmac_16x16x32kernel( data_type *A, data_type *B, data_type *C, data_type *D) { int ip = blockIdx.x * blockDim.x + threadIdx.x; __shared__ data_type A_s[16*32]; __shared__ data_type B_s[32*16]; int row = ip % 16; int col = ip / 16; // K=32，每次处理 K_tile=4，所以需要 32/4 = 8 次迭代 int iter = 32 / 4; // 从全局内存加载 A 和 B 到共享内存 for (int i = 0; i < iter; i++) { int sm_offset = i * 64; A_s[ip + sm_offset] = A[ip + sm_offset]; B_s[ip + sm_offset] = B[ip + sm_offset]; } __syncthreads(); // K 维度循环累加 float4_ D_d; for (int i = 0; i < iter; i++) { int sm_offset = i * 64; CLEAR_ACC(D_d); // 注意：每条 mmac 指令前清零 D_d = __builtin_amdgcn_mmac_f32_16x16x4f32( A_s[ip + sm_offset], B_s[ip + sm_offset], D_d); // 累加到全局内存（注意这里是 +=） D[row * 16 + col + 0 * 4] += D_d.x; D[row * 16 + col + 1 * 4] += D_d.y; D[row * 16 + col + 2 * 4] += D_d.z; D[row * 16 + col + 3 * 4] += D_d.w; } } ``` > **注意**：此示例中每次迭代清零再累加是为了演示 K 维度分片的基本模式。更高效的实现中，`CLEAR_ACC` 只调用一次，后续迭代直接累加。 --- ### 5.6 通用 M×N×K 与软件流水线 为了支持任意 K 值并隐藏访存延迟，采用 **双缓冲（double buffering）** 策略——在计算当前 tile 的同时预取下一个 tile： ```cpp __global__ void vmmac_MxNxK_tc_1wrap_kernel( data_type *A, data_type *B, data_type *C, data_type *D, int m, int n, int k) { int ip = blockIdx.x * blockDim.x + threadIdx.x; int row = ip % 16; int col = ip / 16; // ---- 双缓冲共享内存 ---- __shared__ data_type A_s[2][32*4]; __shared__ data_type B_s[2][4*64]; int iter_k = k / 4; // 预加载第一批数据到 buffer[0] #pragma unroll for (int i = 0; i < 2; i++) A_s[0][ip + i*64] = A[(ip % 32) + (ip/32 + 2*i) * 32]; #pragma unroll for (int i = 0; i < 4; i++) B_s[0][ip + i*64] = B[ip + i*64]; __syncthreads(); float4_ D_d[8]; CLEAR_ACC(D_d[0]); /* ... D_d[1..7] */ ; // ---- 主循环：计算 + 预取交替 ---- for (int k_step = 0; k_step < iter_k; k_step++) { int read_idx = k_step % 2; // 当前使用 buffer[read_idx] int write_idx = (k_step + 1) % 2; // 预取到 buffer[write_idx] if (k_step < iter_k - 1) { // 预取下一批数据（与当前计算重叠） #pragma unroll for (int i = 0; i < 2; i++) A_s[write_idx][ip + i*64] = A[(ip % 32) + (k_step+1)*32*4 + (ip/32 + 2*i) * 32]; #pragma unroll for (int i = 0; i < 4; i++) B_s[write_idx][ip + i*64] = B[ip + (k_step+1)*64*4 + i*64]; } // 使用 buffer[read_idx] 发射 8 条 TensorCore 指令 asm volatile("v_mmac_16x16x4_f32 %0, %1, %2, %0\n\t" : "+v"(D_d[0]) : "v"(A_s[read_idx][row + col*32]), "v"(B_s[read_idx][row + col*64])); // ... 其余 7 条指令类似 ... __syncthreads(); } // ---- 写回结果 ---- // D_d[0..3] → 行 0~15；D_d[4..7] → 行 16~31 // 每个 float4_ 覆盖 N 方向 16 列 ... } ``` 通过双缓冲，全局内存加载与 TensorCore 计算重叠执行，有效隐藏了访存延迟。 --- ### 5.7 双精度 TensorCore DCU K100 同样支持双精度（f64）TensorCore，指令为 `v_mmac_16x16x4_f64`。与单精度的关键区别在于必须使用 **内联汇编** 调用，且累加器类型为 `double4_`： ```cpp __device__ void vmmac_32x64x32_microkernel( data_type *A_s, data_type *B_s, data_type *D_s) { double4_ D_d[8]; // ... 清零 8 个累加器 ... double *A_ptr = A_s + row + col * 32; double *B_ptr = B_s + row + col * 64; for (int i = 0; i < 8; i++) { double a0 = A_ptr[0]; double a1 = A_ptr[16]; double b0 = B_ptr[0]; double b1 = B_ptr[16]; double b2 = B_ptr[32]; double b3 = B_ptr[48]; // 双精度内联汇编：4 条指令，每条使用 4 个 B 寄存器 asm volatile( "s_waitcnt lgkmcnt(0)\n\t" "v_mmac_16x16x4_f64 %0, %4, %5, %0\n\t" "v_mmac_16x16x4_f64 %1, %4, %6, %1\n\t" "v_mmac_16x16x4_f64 %2, %4, %7, %2\n\t" "v_mmac_16x16x4_f64 %3, %4, %8, %3\n\t" : "+v"(D_d[0]), "+v"(D_d[1]), "+v"(D_d[2]), "+v"(D_d[3]) : "v"(a0), "v"(b0), "v"(b1), "v"(b2), "v"(b3)); asm volatile( "s_waitcnt lgkmcnt(0)\n\t" "v_mmac_16x16x4_f64 %0, %4, %5, %0\n\t" "v_mmac_16x16x4_f64 %1, %4, %6, %1\n\t" "v_mmac_16x16x4_f64 %2, %4, %7, %2\n\t" "v_mmac_16x16x4_f64 %3, %4, %8, %3\n\t" : "+v"(D_d[4]), "+v"(D_d[5]), "+v"(D_d[6]), "+v"(D_d[7]) : "v"(a1), "v"(b0), "v"(b1), "v"(b2), "v"(b3)); A_ptr += 4 * 32; // 移动到下一个 K tile B_ptr += 4 * 64; } // ... 写回结果 ... } ``` **单精度 vs 双精度对比**： | 特性 | f32 | f64 | | -------------- | -------------------- | ----------------------- | | **指令** | `v_mmac_16x16x4_f32` | `v_mmac_16x16x4_f64` | | **累加器类型** | `float4_` | `double4_` | | **调用方式** | Builtin 或内联汇编 | 仅内联汇编 | | **每周期吞吐** | 更高 | 较低（约为 f32 的 1/2） | --- ### 5.8 TensorCore总结 1. **Block 大小 = 64 的倍数**：与 Warp 大小对齐，充分利用 TensorCore 并行度。 2. **数据布局**：A 用 col-major，B 用 row-major，使共享内存加载为 coalesced access。 3. **多累加器**：根据 M 和 N 扩展累加器数量（数量 = `ceil(M/16) × ceil(N/16)`），提升指令级并行度。 4. **双缓冲**：使用两套共享内存交替读写，将全局内存加载与 TensorCore 计算重叠。 5. **优先 Builtin，按需切汇编**：f32 场景用 Builtin；f64 或需要精细控制寄存器时用内联汇编。 6. **`s_waitcnt` 保序**：在发射第一条 mmac 指令前用 `s_waitcnt lgkmcnt(0)` 确保 LDS/全局/常量内存的异步读取全部完成（详见 5.3 节说明），避免读到未就绪的脏数据。 --- ## 6. 总结与优化路线图 本文从硬件配置、拓扑分析、带宽测试、性能剖析到 TensorCore 编程，完整走了一遍 DCU K100 节点的上手流程。核心结论与优化建议汇总如下： ### 硬件与拓扑 | 要点 | 关键数据 | | --------------------- | -------------------------------- | | CPU | Hygon C86 7381 ×2，64 核，8 NUMA | | DCU | K100 ×8，架构 gfx926 | | 组内带宽（单向/双向） | 28 / 55 GB/s | | 跨组带宽（单向/双向） | 11 / 22 GB/s | | PCIe 利用率 | 87%~89% | ### 优化路线图 1. **NUMA 亲和性**：确保进程与 DCU 在同一 NUMA 节点内，消除跨节点通信瓶颈。 2. **访存优化**：针对 ALU 利用率低的 Kernel（如文中的描述符导数计算），优先改善访存模式，利用共享内存（LDS）和合并访存。 3. **循环展开**：增加指令级并行度，配合寄存器压力监控，在 vgpr 不溢出的前提下适度展开。 4. **TensorCore 引入**：以 64 的倍数为粒度重构 Kernel 计算框架，按"标准 Kernel → 单条 TC → K 展开 → 多累加器 → 双缓冲软件流水线"的路线逐步优化。 5. **计算-通信重叠**：利用多 Stream 机制，将 PCIe 数据传输与 Kernel 执行流水线化，隐藏通信延迟。 --- ## 参考文献 1. AMD. *AMD ROCm 软件* [EB/OL]. [2025-11-20]. https://www.amd.com/zh-cn/products/software/rocm.html 2. AMD. *Using ROCm Bandwidth Test — ROCm Bandwidth Test Documentation* [EB/OL]. [2025-11-20]. https://rocm.docs.amd.com/projects/rocm_bandwidth_test/en/master/how-to/using_rocm_bandwidth_test.html 3. CUI Chen, WU Di, TAO Yerong, 等. *多 GPU 系统高速互联技术与拓扑结构发展现状研究* [J]. Aero Weaponry（航空兵器）, 2024, 31(1): 23-31. --- *本文基于 DCU K100 节点的实际测试数据整理而成，数据和结论仅供参考。不同硬件批次和驱动版本可能导致性能差异，建议以实际测量结果为准。*