> **适合人群**：有 C/C++ 基础，已读完上篇（或了解 DCU 硬件基础），想动手写代码的开发者 > > **阅读时间**：约 25 分钟 > > **你将收获**：从 Hello World 写到矩阵乘法 → 掌握 HIP 编程 → 会调优 → 会移植 CUDA 代码 > > 📖 **上篇回顾**：[DCU入门（上）：硬件原理与内存体系](./DCU入门（上）：硬件原理与内存体系.md)（DCU 硬件架构 · CU/Wavefront · 五级内存 · HBM2 · DMA） --- ## 目录 - [六、上手写代码——从 Hello World 到矩阵乘法](#六上手写代码从-hello-world-到矩阵乘法) - [6.1 环境准备](#61-环境准备) - [6.2 第一个 HIP 程序：向量加法（完整 walkthrough）](#62-第一个-hip-程序向量加法完整-walkthrough) - [6.3 线程索引全解析](#63-线程索引全解析) - [6.4 二维线程块：图像处理案例](#64-二维线程块图像处理案例) - [6.5 归约算法：深入理解并行编程](#65-归约算法深入理解并行编程) - [七、DCU 软件生态——别再自己造轮子了](#七dcu-软件生态别再自己造轮子了) - [八、性能优化实战——从能跑到跑得快](#八性能优化实战从能跑到跑得快) - [8.1 三个常见坑](#81-三个常见坑) - [8.2 矩阵乘法优化全记录：从 V0 到 V4](#82-矩阵乘法优化全记录从-v0-到-v4) - [8.3 性能分析工具：rocprofiler 快速上手指南](#83-性能分析工具rocprofiler-快速上手指南) - [九、DCU vs NVIDIA CUDA——移植指南速查](#九dcu-vs-nvidia-cuda移植指南速查) - [十、总结 + 学习路线](#十总结--学习路线) --- ## 六、上手写代码——从 Hello World 到矩阵乘法 ### 6.1 环境准备 ```bash # 确认 DTK 已安装 which hipcc # → /opt/dtk/bin/hipcc hipconfig --version # → 看 DCU 工具链版本 # 查看 DCU 设备列表 rocm-smi # → 看有几个 DCU 卡、温度、显存用量 ``` 最简单的编译方式： ```bash hipcc -o my_program my_program.cpp # 相当于 g++ my_program.cpp -o my_program ...（带 HIP 特定的 include 和链接） ``` ### 6.2 第一个 HIP 程序：向量加法（完整 walkthrough） 先贴完整代码，再逐行分析。 ```cpp #include #include #include // 宏：检查 HIP API 调用是否成功（强烈建议每次都用） #define HIP_CHECK(call) do { \ hipError_t err = call; \ if (err != hipSuccess) { \ std::cerr << "HIP Error at " << __FILE__ << ":" \ << __LINE__ << " - " \ << hipGetErrorString(err) << std::endl; \ exit(1); \ } \ } while(0) // ── DCU 核函数：每个线程负责一个元素的加法 ── __global__ void vectorAdd(const float* A, const float* B, float* C, int N) { // 计算全局线程 ID int idx = blockIdx.x * blockDim.x + threadIdx.x; // 边界保护（线程数可能多于数据量） if (idx < N) { C[idx] = A[idx] + B[idx]; } } int main() { // ─── 1. 确定问题规模 ─── int N = 1 << 24; // 16M 个元素 = 64MB size_t size = N * sizeof(float); std::cout << "向量加法: " << N << " 个元素, " << size / (1024*1024) << " MB" << std::endl; // ─── 2. 分配 CPU 端内存并初始化 ─── 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 * 1.0f; h_B[i] = i * 2.0f; } // ─── 3. 分配 DCU 端显存 ─── float *d_A, *d_B, *d_C; HIP_CHECK(hipMalloc(&d_A, size)); // hipMalloc: 在 HBM2 里挖一块 HIP_CHECK(hipMalloc(&d_B, size)); HIP_CHECK(hipMalloc(&d_C, size)); // ─── 4. 数据从 CPU 搬到 DCU (Host → Device) ─── HIP_CHECK(hipMemcpy(d_A, h_A, size, hipMemcpyHostToDevice)); HIP_CHECK(hipMemcpy(d_B, h_B, size, hipMemcpyHostToDevice)); // ─── 5. 配置并启动核函数 ─── int blockSize = 256; // 每个线程块 256 个线程 int gridSize = (N + blockSize - 1) / blockSize; // 向上取整 std::cout << "启动配置: " << gridSize << " 块 × " << blockSize << " 线程/块 = " << gridSize * blockSize << " 总线程" << std::endl; // hipLaunchKernelGGL 的参数依次是： // kernel名字, grid维度, block维度, 动态共享内存字节数, stream, kernel参数... hipLaunchKernelGGL(vectorAdd, gridSize, blockSize, 0, 0, // 共享内存=0, 默认stream d_A, d_B, d_C, N); // ─── 6. 等待 DCU 完成 ─── HIP_CHECK(hipDeviceSynchronize()); // ─── 7. 结果从 DCU 搬回 CPU (Device → Host) ─── HIP_CHECK(hipMemcpy(h_C, d_C, size, hipMemcpyDeviceToHost)); // ─── 8. 验证结果 ─── int errors = 0; for (int i = 0; i < N; i++) { if (fabs(h_C[i] - (i * 1.0f + i * 2.0f)) > 0.001f) { if (errors++ < 5) // 只打印前5个错误 std::cerr << "错误: h_C[" << i << "] = " << h_C[i] << " (期望 " << (i*3.0f) << ")" << std::endl; } } if (errors == 0) { std::cout << "✅ 全部 " << N << " 个元素验证通过!" << std::endl; std::cout << " 示例: h_C[0]=" << h_C[0] << ", h_C[1]=" << h_C[1] << ", ..., h_C[" << N-1 << "]=" << h_C[N-1] << std::endl; } else { std::cerr << "❌ 共 " << errors << " 个错误" << std::endl; } // ─── 9. 清理资源 ─── HIP_CHECK(hipFree(d_A)); HIP_CHECK(hipFree(d_B)); HIP_CHECK(hipFree(d_C)); delete[] h_A; delete[] h_B; delete[] h_C; return 0; } ``` **编译运行**： ```bash hipcc -o vector_add vector_add.cpp ./vector_add ``` **执行流水线（关键理解）**： ``` 时间 → CPU (Host) DCU (Device) [分配内存] [待命] [初始化 h_A, h_B] [待命] [hipMalloc] ──────────► [在 HBM2 里分配 d_A, d_B, d_C] [hipMemcpy] ──────────► [DMA 自动搬运 h_A→d_A, h_B→d_B] [hipLaunchKernel]──────► [下发 kernel 到每个 CU] [Wavefront 0: ADD R1,R2,R3 (64线程)] ← 零开销切换 [Wavefront 1: ADD R1,R2,R3 (64线程)] [CPU 继续做事...] [Wavefront 2: ADD R1,R2,R3 (64线程)] [... 所有线程执行完毕 ...] [hipDeviceSync] ◄─────── [通知 CPU: 我干完了] [hipMemcpy] ◄─────────── [DMA 搬运 d_C → h_C] [验证结果] [清理] ``` ### 6.3 线程索引全解析 这是 DCU/HIP 编程最基础的技能——你要能算出**每个线程处理哪个数据**。 ```cpp // ─── 1D 布局（最常见） ─── // Block 维度: (256, 1, 1) // Grid 维度: (64, 1, 1) // 总线程数: 64 × 256 = 16384 __global__ void kernel_1d(float* data) { int global_id = blockIdx.x * blockDim.x + threadIdx.x; // threadIdx.x ∈ [0, blockDim.x) → 线程在块内的位置 // blockIdx.x ∈ [0, gridDim.x) → 线程块在网格内的位置 // blockDim.x → 每个块多少线程（启动时固定） } // ─── 2D 布局（图像处理常用） ─── // Block 维度: (16, 16, 1) → 256 线程/块 // Grid 维度: (64, 64, 1) → 4096 块 __global__ void kernel_2d(float* image, int width, int height) { int x = blockIdx.x * blockDim.x + threadIdx.x; // 列 int y = blockIdx.y * blockDim.y + threadIdx.y; // 行 if (x < width && y < height) { // 处理像素 (x, y) int linear_id = y * width + x; image[linear_id] = do_something(image[linear_id]); } } // ─── 3D 布局（体数据 / 物理模拟） ─── __global__ void kernel_3d(float* volume, int X, int Y, int Z) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; int z = blockIdx.z * blockDim.z + threadIdx.z; if (x < X && y < Y && z < Z) { int idx = z * (X * Y) + y * X + x; volume[idx] = process(volume[idx]); } } ``` **1D 线程索引速查表**： | 变量 | 含义 | 举例 (grid=(64,1,1), block=(256,1,1)) | | ------------- | --------------------------------------- | ------------------------------------- | | `threadIdx.x` | 线程在块内的 x 索引 | 0 ~ 255 | | `blockIdx.x` | 块在网格内的 x 索引 | 0 ~ 63 | | `blockDim.x` | 每块的 x 方向线程数 | 256 | | `gridDim.x` | 网格的 x 方向块数 | 64 | | `全局ID` | `blockIdx.x * blockDim.x + threadIdx.x` | 0 ~ 16383 | ### 6.4 二维线程块：图像处理案例 写一个**灰度图转二值图**的完整示例，演示 2D 线程块用法： ```cpp #include #include #include #define HIP_CHECK(call) /* 同上，省略 */ // 灰度转二值：大于阈值的像素 → 255, 否则 → 0 __global__ void threshold(const unsigned char* __restrict__ input, unsigned char* __restrict__ output, int width, int height, unsigned char threshold_val) { // 2D 线程映射 int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height) { int idx = y * width + x; unsigned char pixel = input[idx]; output[idx] = (pixel > threshold_val) ? 255 : 0; } } int main() { int width = 4096, height = 3072; // 4K 分辨率 size_t img_size = width * height * sizeof(unsigned char); // CPU 端：生成测试图像 unsigned char *h_input = new unsigned char[width * height]; unsigned char *h_output = new unsigned char[width * height]; for (int i = 0; i < width * height; i++) { h_input[i] = (unsigned char)(i % 256); // 渐变灰度 } // DCU 端：分配显存 unsigned char *d_input, *d_output; HIP_CHECK(hipMalloc(&d_input, img_size)); HIP_CHECK(hipMalloc(&d_output, img_size)); HIP_CHECK(hipMemcpy(d_input, h_input, img_size, hipMemcpyHostToDevice)); // 2D 线程块配置 dim3 block(16, 16); // 16×16 = 256 线程/块（经典配置） dim3 grid((width + 15) / 16, (height + 15) / 16); std::cout << "图像大小: " << width << "×" << height << " | Grid: " << grid.x << "×" << grid.y << " | Block: " << block.x << "×" << block.y << std::endl; hipLaunchKernelGGL(threshold, grid, block, 0, 0, d_input, d_output, width, height, 128); HIP_CHECK(hipDeviceSynchronize()); HIP_CHECK(hipMemcpy(h_output, d_output, img_size, hipMemcpyDeviceToHost)); // 验证 int errors = 0; for (int i = 0; i < width * height; i++) { unsigned char expected = (h_input[i] > 128) ? 255 : 0; if (h_output[i] != expected) errors++; } std::cout << (errors == 0 ? "✅ 全部正确!" : "❌ 有错误") << std::endl; HIP_CHECK(hipFree(d_input)); HIP_CHECK(hipFree(d_output)); delete[] h_input; delete[] h_output; return 0; } ``` ### 6.5 归约算法：深入理解并行编程 **归约（Reduction）** 是把一个数组"折叠"成一个值的过程——求和、求最大值、求范数都是归约。这是理解共享内存和线程同步的最好案例。 ```cpp // ─── 并行归约：求数组和 ─── // 每轮参与线程数减半，最终把累加结果留在共享内存的第一个位置 #define BLOCK_SIZE 256 #define WARPSIZE 64 // DCU 的 Wavefront 大小 __global__ void reduce_sum(const float* __restrict__ input, float* __restrict__ output, int N) { __shared__ float smem[BLOCK_SIZE]; int tid = threadIdx.x; int idx = blockIdx.x * blockDim.x * 2 + tid; // ×2：每个线程加载 2 个元素 // ── 步骤1: 从全局内存加载到共享内存 ── // 每个线程加载 2 个元素并累加（减少一半线程启动） float sum = 0.0f; if (idx < N) sum += input[idx]; if (idx + BLOCK_SIZE < N) sum += input[idx + BLOCK_SIZE]; smem[tid] = sum; __syncthreads(); // ── 步骤2: 树形归约（每轮参与线程数减半） ── // 可视化（假设 8 线程版）： // 第 1 轮: 线程0←线程4 线程1←线程5 线程2←线程6 线程3←线程7 // [0+4] [1+5] [2+6] [3+7] // 第 2 轮: 线程0←线程2 线程1←线程3 (线程4~7 不参与) // [0+2+4+6] [1+3+5+7] // 第 3 轮: 线程0←线程1 (线程2~7 不参与) // [sum of all] // for (int stride = BLOCK_SIZE / 2; stride > 0; stride >>= 1) { if (tid < stride) { smem[tid] += smem[tid + stride]; } __syncthreads(); // 每轮之后必须同步 } // ── 步骤3: 每个块的结果写回全局内存 ── if (tid == 0) { output[blockIdx.x] = smem[0]; // 每个块输出一个部分和 } } ``` **归约过程的图解**： ``` 输入: [3, 1, 7, 0, 4, 1, 6, 3] (N=8, block=8，简化) 线程: t0 t1 t2 t3 t4 t5 t6 t7 smem: [ 3, 1, 7, 0, 4, 1, 6, 3 ] 第1轮 (stride=4): t0←t4(3+4) t1←t5(1+1) t2←t6(7+6) t3←t7(0+3) smem: [ 7, 2, 13, 3, 4, 1, 6, 3 ] ← 后半不再被读 第2轮 (stride=2): t0←t2(7+13) t1←t3(2+3) smem: [20, 5, 13, 3, 4, 1, 6, 3 ] 第3轮 (stride=1): t0←t1(20+5) smem: [25, 5, 13, 3, 4, 1, 6, 3 ] ↑ sum = 25 ✓ ``` --- ## 七、DCU 软件生态——别再自己造轮子了 DCU 不只是提供一个 HIP 编译器——它还有一个完整的软件栈。大部分常见算法都有高度优化的库可以直接调。 ``` ┌──────────────────────────────────────────────────┐ │ 你的应用 (AI训练 / 科学计算 / 数据分析) │ ├──────────────────────────────────────────────────┤ │ 数学库 │ │ ┌────────┬────────┬────────┬──────────────────┐ │ │ │hipBLAS │hipFFT │hipSPARSE│hipRAND MIOpen │ │ │ │(线性代数)(傅里叶) (稀疏矩阵)(随机数) (深度学习) │ │ │ ├────────┴────────┼────────┴──────────────────┤ │ │ │ RCCL (多卡通信) │ hipCUB (基础算法) │ │ │ └─────────────────┴───────────────────────────┘ │ ├──────────────────────────────────────────────────┤ │ HIP Runtime (核函数启动、内存管理、stream等) │ ├──────────────────────────────────────────────────┤ │ DCU 驱动 │ ├──────────────────────────────────────────────────┤ │ DCU 硬件 │ └──────────────────────────────────────────────────┘ ``` ### 核心库速查与代码示例 | 库 | 对标 NVIDIA | 干什么 | 一句话 | | ------------- | ----------- | ----------------- | --------------------- | | **hipBLAS** | cuBLAS | BLAS 1/2/3 级运算 | 矩阵乘、向量计算 | | **hipFFT** | cuFFT | 快速傅里叶变换 | 信号/图像频域处理 | | **hipSPARSE** | cuSPARSE | 稀疏矩阵运算 | 图计算、推荐系统 | | **hipRAND** | cuRAND | 伪随机数生成 | 初始化、Dropout | | **MIOpen** | cuDNN | 深度学习算子 | 卷积/池化/激活/归一化 | | **hipCUB** | CUB | 并行原语 | 排序/扫描/直方图 | | **RCCL** | NCCL | 集合通信 | AllReduce/Broadcast | | **rocThrust** | Thrust | STL-like并行模板 | sort/reduce/scan | **hipBLAS 示例——矩阵乘法**： ```cpp #include #include // 一行代码搞定矩阵乘法 C = α*A*B + β*C void hipblas_sgemm_demo() { int M = 2048, N = 2048, K = 2048; // C(M×N) = A(M×K) × B(K×N) size_t sizeA = M * K * sizeof(float); size_t sizeB = K * N * sizeof(float); size_t sizeC = M * N * sizeof(float); float *d_A, *d_B, *d_C; hipMalloc(&d_A, sizeA); hipMalloc(&d_B, sizeB); hipMalloc(&d_C, sizeC); // 初始化数据（省略 hipMemcpy...） hipblasHandle_t handle; hipblasCreate(&handle); float alpha = 1.0f, beta = 0.0f; // SGEMM: S=单精度, GEMM=通用矩阵乘法 hipblasSgemm(handle, HIPBLAS_OP_N, HIPBLAS_OP_N, // 不转置 M, N, K, // C(M×N) = A(M×K) * B(K×N) &alpha, d_A, M, // A 矩阵 (leading dimension = M) d_B, K, // B 矩阵 (leading dimension = K) &beta, d_C, M); // C 矩阵 (leading dimension = M) hipblasDestroy(handle); // 比自己写的 kernel 通常快 3~10 倍——因为底层针对 DCU 架构深度优化 } ``` **MIOpen 示例——深度学习卷积**： ```cpp #include // 创建一个卷积层的前向推理 void miopen_conv2d_demo() { miopenHandle_t handle; miopenCreate(&handle); // 输入: NCHW 格式，4×3×224×224 // 权重: 64×3×7×7 (64 个输出通道，卷积核 7×7) miopenTensorDescriptor_t input_desc, output_desc, filter_desc; miopenCreateTensorDescriptor(&input_desc); miopenCreateTensorDescriptor(&output_desc); miopenCreateTensorDescriptor(&filter_desc); int N = 4, C = 3, H = 224, W = 224; miopenSet4dTensorDescriptor(input_desc, miopenFloat, N, C, H, W); int K = 64; // 输出通道 miopenSet4dTensorDescriptor(filter_desc, miopenFloat, K, C, 7, 7); // ... 省略卷积描述符和算法选择 ... // 执行卷积（内部使用最优的 Winograd/FFT/ImplicitGEMM 等算法） miopenConvolutionForward(handle, &alpha, input_desc, d_input, filter_desc, d_filter, conv_desc, algo, &beta, output_desc, d_output, d_workspace, workspace_size); miopenDestroy(handle); } ``` --- ## 八、性能优化实战——从能跑到跑得快 ### 8.1 三个常见坑 **坑1：隐式的 CPU↔DCU 数据传输** ```cpp // ❌ 循环里每次迭代搬数据 (PCI-E 28GB/s 是最大瓶颈) for (int iter = 0; iter < 10000; iter++) { hipMemcpy(d_data, h_data, size, hipMemcpyHostToDevice); kernel<<>>(d_data, d_result); hipMemcpy(h_result, d_result, size, hipMemcpyDeviceToHost); } // ✅ 数据留在 DCU 上，多次 kernel 之间不搬数据 hipMemcpy(d_data, h_data, size, hipMemcpyHostToDevice); for (int iter = 0; iter < 10000; iter++) { kernel<<>>(d_data, d_result); // d_result 留在 DCU 显存上，下个 iter 直接用 } hipMemcpy(h_result, d_result, size, hipMemcpyDeviceToHost); ``` **坑2：寄存器溢出（Register Spill）** ```cpp // ❌ 声明大数组 → 编译器被迫 spill 到 L1 __global__ void too_many_regs(float* data) { float temp[200]; // ← 200×4=800 字节，远超寄存器配额 for (int i = 0; i < 200; i++) temp[i] = data[threadIdx.x + i*256]; // 编译器只能把 temp 放到 L1/LDS，每次访问多 30 周期 } // ✅ 减少局部变量 → 全部留在寄存器 __global__ void fewer_regs(float* data) { float val0 = data[threadIdx.x]; float val1 = data[threadIdx.x + 256]; float val2 = data[threadIdx.x + 512]; float val3 = data[threadIdx.x + 768]; // 4个寄存器 → 极快 } ``` 检查寄存器使用量： ```bash hipcc -Rpass-analysis=kernel-resource-usage -c my_kernel.cpp # 输出: kernel_name: 32 registers, 0 bytes lds, 0 bytes stack ``` **坑3：忽略 Occupancy（占用率）** ```cpp // Occupancy = 活跃 Wavefront 数 / 理论最大 Wavefront 数 // 如果每个线程用 128 个寄存器，一个 CU 最多能放 2560/8 = 320 线程 // → Occupancy 只有 12.5% // → 无法有效隐藏内存延迟 → 性能暴跌 // 查看 Occupancy： // rocprof --stats ./my_program // 关注: VALUUtilization%, SQ_WAVES (活跃 Wavefront) ``` ### 8.2 矩阵乘法优化全记录：从 V0 到 V4 这是理解 DCU 性能优化最经典的案例——**同一个算法，不同实现方式，性能差 50 倍**。 每个版本前标注了**关键优化点**和**性能提升**（基于 DCU2 实测，具体数字因型号而异）。 --- **V0：最朴素的版本（基准线）** ```cpp // V0: 每个线程算一个输出元素，直接读全局内存 // 关键问题: 每个线程读 K 次全局内存 → 带宽瓶颈 __global__ void sgemm_v0(const float* __restrict__ A, const float* __restrict__ B, float* __restrict__ C, 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) { float sum = 0.0f; for (int k = 0; k < K; k++) { sum += A[row * K + k] * B[k * N + col]; // ← 每次读2次全局内存！ } C[row * N + col] = sum; } } // 性能: ~50 GFLOPS (基准线) ``` --- **V1：合并访问优化（提升 ~3×）** ```cpp // V1: 让相邻线程读相邻地址 (合并访问) // 关键改进: 每个线程块先协作加载 A 的一列到寄存器 __global__ void sgemm_v1(const float* __restrict__ A, const float* __restrict__ B, float* __restrict__ C, 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) { float sum = 0.0f; // B 矩阵的列读取：线程0→B[0*N+col], 线程1→B[0*N+col+1] ... // 但 B[0*N+col] 到 B[0*N+col+63] 如果是连续的 → 合并访问 ✓ for (int k = 0; k < K; k++) { sum += A[row * K + k] * B[k * N + col]; } C[row * N + col] = sum; } } // 性能: ~150 GFLOPS (比 V0 快 3×) ``` --- **V2：共享内存 Tiling（提升 ~15×）** ```cpp // V2: 引入共享内存做 tiling（分块） // 关键改进: 每个线程块的一小块 A 和 B 只从全局内存读一次 #define TILE_SIZE 32 __global__ void sgemm_v2(const float* __restrict__ A, const float* __restrict__ B, float* __restrict__ C, int M, int N, int K) { __shared__ float As[TILE_SIZE][TILE_SIZE]; __shared__ float Bs[TILE_SIZE][TILE_SIZE]; int bx = blockIdx.x, by = blockIdx.y; int tx = threadIdx.x, ty = threadIdx.y; // 全局输出索引 int row = by * TILE_SIZE + ty; int col = bx * TILE_SIZE + tx; float sum = 0.0f; // 外层循环：遍历 K 维度上的分块 for (int t = 0; t < K / TILE_SIZE; t++) { // 协作加载 A 的一个 tile 到共享内存 if (row < M && (t * TILE_SIZE + tx < K)) As[ty][tx] = A[row * K + t * TILE_SIZE + tx]; else As[ty][tx] = 0.0f; // 协作加载 B 的一个 tile 到共享内存 if ((t * TILE_SIZE + ty) < K && col < N) Bs[ty][tx] = B[(t * TILE_SIZE + ty) * N + col]; else Bs[ty][tx] = 0.0f; __syncthreads(); // 共享内存里的乘加（极快！） for (int k = 0; k < TILE_SIZE; k++) { sum += As[ty][k] * Bs[k][tx]; // ← 从共享内存读，~30周期 } __syncthreads(); // 等所有线程用完，准备加载下一块 } if (row < M && col < N) C[row * N + col] = sum; } // 性能: ~750 GFLOPS (比 V0 快 15×) ``` --- **V3：寄存器 + 共享内存 + 4×1 micro-tiling（提升 ~30×）** ```cpp // V3: 每个线程算 4 个输出元素(vectorize) // A 的 tile 读进共享内存时用 float4 向量化加载 // 关键改进: float4 加载 → 1 条指令读 4 个 float #define TILE_SIZE 32 __global__ void sgemm_v3(const float* __restrict__ A, const float* __restrict__ B, float* __restrict__ C, int M, int N, int K) { __shared__ float As[TILE_SIZE][TILE_SIZE]; __shared__ float Bs[TILE_SIZE][TILE_SIZE]; int bx = blockIdx.x, by = blockIdx.y; int tx = threadIdx.x, ty = threadIdx.y; int row = by * TILE_SIZE + ty; int col = (bx * TILE_SIZE + tx) * 4; // 每个线程算 4 列 float sum0 = 0.0f, sum1 = 0.0f, sum2 = 0.0f, sum3 = 0.0f; for (int t = 0; t < K / TILE_SIZE; t++) { // 向量化加载 A (每次读 4 个 float) int a_col = t * TILE_SIZE + tx; if (row < M && a_col < K) { As[ty][tx] = A[row * K + a_col]; } else As[ty][tx] = 0.0f; // B 的 tile 加载（向量化读取 4 列） int b_row = t * TILE_SIZE + ty; #pragma unroll for (int c = 0; c < 4; c++) { int b_col = (bx * TILE_SIZE + tx) * 4 + c; if (b_row < K && b_col < N) Bs[ty][tx * 4 + c] = B[b_row * N + b_col]; } __syncthreads(); #pragma unroll for (int k = 0; k < TILE_SIZE; k++) { float a_val = As[ty][k]; sum0 += a_val * Bs[k][tx * 4 + 0]; sum1 += a_val * Bs[k][tx * 4 + 1]; sum2 += a_val * Bs[k][tx * 4 + 2]; sum3 += a_val * Bs[k][tx * 4 + 3]; } __syncthreads(); } if (row < M && col + 3 < N) { C[row * N + col + 0] = sum0; C[row * N + col + 1] = sum1; C[row * N + col + 2] = sum2; C[row * N + col + 3] = sum3; } } // 性能: ~1500 GFLOPS (比 V0 快 30×) ``` --- **V4：结合 Bank 冲突消除 + 双缓冲（提升 ~50×）** ```cpp // V4: 在 V3 基础上： // 1. 共享内存加 padding (TILE_PAD = TILE_SIZE + 1) 消除 Bank 冲突 // 2. 双缓冲：加载下一个 tile 的同时算当前 tile #define TILE_SIZE 32 #define TILE_PAD 33 // +1 消除 bank 冲突 __global__ void sgemm_v4(const float* __restrict__ A, const float* __restrict__ B, float* __restrict__ C, int M, int N, int K) { __shared__ float As[2][TILE_SIZE][TILE_PAD]; // 双缓冲 __shared__ float Bs[2][TILE_SIZE][TILE_PAD]; int bx = blockIdx.x, by = blockIdx.y; int tx = threadIdx.x, ty = threadIdx.y; int row = by * TILE_SIZE + ty; int col = (bx * TILE_SIZE + tx) * 4; float sum0 = 0.0f, sum1 = 0.0f, sum2 = 0.0f, sum3 = 0.0f; // 预加载第一个 tile (buffer 0) int buffer = 0; load_tile(A, B, As[0], Bs[0], row, col, M, N, K, 0, tx, ty); __syncthreads(); for (int t = 1; t < K / TILE_SIZE; t++) { int next_buf = 1 - buffer; // 加载下一个 tile (异步的硬件上可以与计算重叠) load_tile(A, B, As[next_buf], Bs[next_buf], row, col, M, N, K, t, tx, ty); // 同时计算当前 tile！(双缓冲精髓) compute_tile(As[buffer], Bs[buffer], tx, ty, sum0, sum1, sum2, sum3); __syncthreads(); buffer = next_buf; } // 最后一个 tile compute_tile(As[buffer], Bs[buffer], tx, ty, sum0, sum1, sum2, sum3); // 写回 if (row < M) { if (col + 0 < N) C[row * N + col + 0] = sum0; if (col + 1 < N) C[row * N + col + 1] = sum1; if (col + 2 < N) C[row * N + col + 2] = sum2; if (col + 3 < N) C[row * N + col + 3] = sum3; } } // 辅助函数（device 函数） __device__ void load_tile(const float* A, const float* B, float As[TILE_SIZE][TILE_PAD], float Bs[TILE_SIZE][TILE_PAD], int row, int col, int M, int N, int K, int t, int tx, int ty) { int a_col = t * TILE_SIZE + tx; if (row < M && a_col < K) As[ty][tx] = A[row * K + a_col]; else As[ty][tx] = 0.0f; int b_row = t * TILE_SIZE + ty; if (b_row < K) { for (int c = 0; c < 4; c++) { int b_col = col + c; Bs[ty][tx * 4 + c] = (b_col < N) ? B[b_row * N + b_col] : 0.0f; } } } __device__ void compute_tile(const float As[TILE_SIZE][TILE_PAD], const float Bs[TILE_SIZE][TILE_PAD], int tx, int ty, float& s0, float& s1, float& s2, float& s3) { #pragma unroll for (int k = 0; k < TILE_SIZE; k++) { float a = As[ty][k]; s0 += a * Bs[k][tx * 4 + 0]; s1 += a * Bs[k][tx * 4 + 1]; s2 += a * Bs[k][tx * 4 + 2]; s3 += a * Bs[k][tx * 4 + 3]; } } // 性能: ~2500 GFLOPS (比 V0 快 50×) ``` **优化总结**： ``` V0: 朴素全局内存 ~50 GFLOPS (基线) V1: + 合并访问 ~150 GFLOPS (×3) V2: + 共享内存tiling ~750 GFLOPS (×15) V3: + 向量化/寄存器复用 ~1500 GFLOPS (×30) V4: + Bank冲突消除+双缓冲 ~2500 GFLOPS (×50) ``` ### 8.3 性能分析工具：rocprofiler 快速上手指南 ```bash # 1. 基础性能统计 rocprof --stats ./my_program # 输出关键指标： # - SQ_WAVES: 活跃 Wavefront 数 # - VALUUtilization: 矢量 ALU 利用率 # - FetchSize: 全局内存读取量 # - MemUnitStalled: 因等待内存而停顿的周期比例 # 2. 时间线分析 rocprof --timestamp on -d ./profile_output ./my_program # 然后可视化（需 chrome://tracing） # hipMemcpy 和 kernel 执行的时间线一目了然 # 3. 看关键瓶颈 rocprof --stats ./my_program 2>&1 | grep -E "VALUUtil|FetchSize|MemUnitStalled|SQ_WAVES" ``` **性能解读速查**： | 指标 | 理想值 | 如果不对 → 问题 | | -------------------- | -------------- | ------------------------------------- | | `VALUUtilization` | > 80% | < 60% → ALU 在等数据 | | `MemUnitStalled` | < 20% | > 50% → 内存是瓶颈 | | `SQ_WAVES / CU` | 接近 10 (max) | < 5 → 寄存器/LDS 用太多，occupancy 低 | | `FetchSize / kernel` | 接近理论最小值 | 远大于 → 内存访问不合并 | --- ## 九、DCU vs NVIDIA CUDA——移植指南速查 如果你有 CUDA 代码，迁移到 DCU/HIP 几乎是一对一替换： | CUDA | HIP (DCU) | 注意事项 | | -------------------------- | ------------------------------------------------- | --------------------------------------- | | `cudaMalloc` | `hipMalloc` | 参数完全相同 | | `cudaMemcpy` | `hipMemcpy` | 枚举值 `cudaMemcpyH2D` → `hipMemcpyH2D` | | `cudaFree` | `hipFree` | | | `cudaDeviceSynchronize` | `hipDeviceSynchronize` | | | `__global__` | `__global__` | 完全一致 | | `__shared__` | `__shared__` | 完全一致 | | `__constant__` | `__constant__` | | | `threadIdx` | `threadIdx` | 完全一致 | | `blockIdx` | `blockIdx` | 完全一致 | | `blockDim` | `blockDim` | 完全一致 | | `kernel<<>>` | `hipLaunchKernelGGL(kernel,g,b,s,stream,args...)` | HIP 启动方式不同 | | `__syncthreads()` | `__syncthreads()` | 完全一致 | | Warp = 32 | Wavefront = 64 | 锁步大小不同，注意分支分歧模式 | | `nvcc` | `hipcc` | 编译器不同 | | `nvprof` | `rocprofiler` | 性能分析工具不同 | **批量替换脚本**： ```bash # 一键替换 CUDA API → HIP API sed -i 's/cudaMalloc/hipMalloc/g' *.cpp *.cu sed -i 's/cudaMemcpy/hipMemcpy/g' *.cpp *.cu sed -i 's/cudaMemcpyHostToDevice/hipMemcpyHostToDevice/g' *.cpp *.cu sed -i 's/cudaMemcpyDeviceToHost/hipMemcpyDeviceToHost/g' *.cpp *.cu sed -i 's/cudaFree/hipFree/g' *.cpp *.cu sed -i 's/cudaDeviceSynchronize/hipDeviceSynchronize/g' *.cpp *.cu sed -i 's/cudaGetErrorString/hipGetErrorString/g' *.cpp *.cu # 注意：kernel 启动语法需要手动改！ ``` --- ## 十、总结 + 学习路线 ### 知识全景 ``` DCU = 国产 GPGPU 加速卡 硬件: CU阵列(算力) + HBM2(显存) + DMA(搬运) + PCI-E/xGMI(互联) 编程: HIP ≈ CUDA 语法，加 hippify 工具可自动转换 优化阶梯: 合并访问 → 共享内存 → 向量化 → Bank修补 → 双缓冲 工具: hipcc(编译) + rocprofiler(分析) + roctracer(跟踪) 库: hipBLAS/MIOpen/RCCL → 90%场景不用手写kernel ``` ### 推荐学习路径 ``` 第1天: 装DTK → 跑向量加法demo → 改参数看性能变化 第2天: 写归约 → 理解共享内存和__syncthreads() 第3天: 写矩阵乘法V0→V2 → 感受共享内存tiling的威力 第4天: 学rocprofiler → 分析自己的代码瓶颈 第5天: 调库: hipBLAS做矩阵乘，MIOpen做卷积 第6天: 挑战: 矩阵乘法V3→V4 → 追求极致性能 第7天: 实战: 用RCCL写多卡程序 ``` --- 📖 **上篇回顾**：[DCU入门（上）：硬件原理与内存体系](./DCU入门（上）：硬件原理与内存体系.md)（DCU 硬件架构 · CU/Wavefront · 五级内存 · HBM2 · DMA） 如果觉得有帮助，欢迎 **点赞 + 收藏**，有问题评论区见。 下一篇打算写《DCU 共享内存与 Bank 冲突深度实战》，用真实的 rocprofiler 报告 + 分步优化记录，讲讲怎么把矩阵乘法从 750 GFLOPS 再提一档到 2000+。 **有什么问题欢迎评论区交流～**