目 录
第 1章 介绍 ………………………………………………………………………
1.1 并行编程 ……………………………………………………………………………… 2
1.2 GPU …………………………………………………………………………………… 3
1.3 ROCm ………………………………………………………………………………… 5
1.4 HIP框架 ……………………………………………………………………………… 6
1.5 本书内容 ……………………………………………………………………………… 6
第 2章 开始使用 HIP编程 ………………………………………………………
2.1 在 HIP中实现“Hello World” ……………………………………………………… 8
2.2 使用 HIP进行数据处理——Vector Add 示例 ……………………………………… 9
2.2.1 并行执行的机会……………………………………………………………… 9
2.2.2 组织线程…………………………………………………………………… 10
2.2.3 使用 HIP API 进行数据传输 ……………………………………………… 12
2.2.4 错误和正确性检查………………………………………………………… 13
2.2.5 综合示例…………………………………………………………………… 14
2.3 总结 ………………………………………………………………………………… 16
第 3章 HIP内核编程………………………………………………………………
3.1 在 HIP内核中调用函数 …………………………………………………………… 18
3.1.1 HIP中的 __global__ 函数 ……………………………………………… 18
3.1.2 HIP中的 __device_ _ 函数 ……………………………………………… 19
3.1.3 HIP中的 __host__ 函数 ………………………………………………… 20
3.1.4 结合 __host__ 和 _ _device__ 函数 ……………………………………… 20
3.2 在 HIP内核中使用模板 …………………………………………………………… 21
3.3 在 HIP中使用结构体 ……………………………………………………………… 22
3.4 总结 ………………………………………………………………………………… 26
第 4章 HIP运行时 API ……………………………………………………………
4.1 HIP内存管理 ……………………………………………………………………… 27
4.1.1 固定内存…………………………………………………………………… 27
4.1.2 统一内存…………………………………………………………………… 29
4.2 HIP流 ……………………………………………………………………………… 30
4.2.1 流的基础知识……………………………………………………………… 31
目 录
4.2.2 基于流的关键 API ………………………………………………………… 31
4.2.3 默认流与非默认流………………………………………………………… 32
4.2.4 并发内核…………………………………………………………………… 33
4.2.5 重叠计算与通信…………………………………………………………… 33
4.3 HIP事件 …………………………………………………………………………… 35
4.3.1 创建 HIP事件 …………………………………………………………… 35
4.3.2 记录 HIP事件 …………………………………………………………… 35
4.3.3 使用 HIP事件计算耗时 ………………………………………………… 36
4.3.4 使用 HIP事件协调操作 ………………………………………………… 36
4.3.5 使用 HIP事件释放内存 ………………………………………………… 36
4.3.6 创建具有特定标志的事件………………………………………………… 37
4.3.7 确保两种处理单元之间的一致性………………………………………… 37
4.3.8 使用 HIP事件的向量加法 ……………………………………………… 37
4.4 总结 ………………………………………………………………………………… 39
第 5章 GPU编程模式 ……………………………………………………………
5.1 二维内核 …………………………………………………………………………… 40
5.2 模板计算 …………………………………………………………………………… 44
5.3 多内核示例——广度优先搜索 …………………………………………………… 47
5.4 CPU-GPU 协同计算——KMeans ………………………………………………… 49
5.5 原子操作——直方图 ……………………………………………………………… 52
5.6 总结 ………………………………………………………………………………… 54
第 6章 AMD GPU内部原理………………………………………………………
6.1 AMD GPU ………………………………………………………………………… 55
6.2 整体架构 …………………………………………………………………………… 57
6.3 命令处理器和 DMA引擎 ………………………………………………………… 58
6.4 工作组分发 ………………………………………………………………………… 58
6.5 指令调度器 ………………………………………………………………………… 60
6.6 SIMD 单元 ………………………………………………………………………… 61
6.7 线程分歧 …………………………………………………………………………… 62
6.8 内存合并 …………………………………………………………………………… 64
目 录
6.9 内存层次结构 ……………………………………………………………………… 65
6.10 AMD RDNA GPUs ……………………………………………………………… 66
6.11 总结 ……………………………………………………………………………… 67
第 7章 HIP工具……………………………………………………………………
7.1 ROCmInfo ………………………………………………………………………… 68
7.2 ROCm SMI ………………………………………………………………………… 71
7.3 ROCm 调试器 ……………………………………………………………………… 73
7.4 ROCm 分析器 ……………………………………………………………………… 76
7.4.1 ROCTracer ………………………………………………………………… 76
7.4.2 ROCProfiler ……………………………………………………………… 78
7.5 ROCm Profiler V2 ………………………………………………………………… 80
7.5.1 应用程序追踪……………………………………………………………… 80
7.5.2 内核分析…………………………………………………………………… 81
7.5.3 ROCSys …………………………………………………………………… 84
7.6 使用 Hipify 将 CUDA程序移植到 HIP…………………………………………… 86
7.6.1 Hipify 工具 ………………………………………………………………… 86
7.6.2 一般 Hipify 指南 ………………………………………………………… 90
7.6.3 矩阵转置的 Hipify ………………………………………………………… 91
7.6.4 常见问题及解决方案……………………………………………………… 93
7.7 总结 ………………………………………………………………………………… 94
第 8章 HIP性能优化………………………………………………………………
8.1 高度并行的工作负载——图像伽马校正 ………………………………………… 95
8.2 固定大小的内核——图像伽马校正 ……………………………………………… 98
8.3 归约——数组求和 ………………………………………………………………… 100
8.4 分块与复用——矩阵乘法 ………………………………………………………… 103
8.5 分块与合并——矩阵转置 ………………………………………………………… 107
8.6 总结 ………………………………………………………………………………… 109
第 9章 ROCm库 ………………………………………………………………
9.1 rocBLAS …………………………………………………………………………… 110
9.1.1 使用 rocBLAS …………………………………………………………… 111
目 录
9.1.2 rocBLAS 函数 …………………………………………………………… 113
9.1.3 异步执行…………………………………………………………………… 114
9.1.4 MI100 上的 rocBLAS …………………………………………………… 115
9.1.5 从传统 BLAS 库迁移 …………………………………………………… 115
9.2 rocSPARSE ………………………………………………………………………… 116
9.2.1 稀疏数据表示……………………………………………………………… 116
9.2.2 rocSPARSE 函数 ………………………………………………………… 118
9.3 rocFFT ……………………………………………………………………………… 120
9.3.1 rocFFT工作流程 ………………………………………………………… 121
9.3.2 FFT执行计划……………………………………………………………… 122
9.4 rocRAND …………………………………………………………………………… 123
9.5 总结 ………………………………………………………………………………… 126
第 10章 多 GPU编程 …………………………………………………………
10.1 HIP设备 API ……………………………………………………………………… 128
10.2 基于流的多 GPU 编程 …………………………………………………………… 130
10.3 基于线程的多 GPU 编程 ………………………………………………………… 132
10.4 基于 MPI 的多 GPU 编程………………………………………………………… 134
10.5 GPU—GPU 通信 ………………………………………………………………… 137
10.6 RCCL ……………………………………………………………………………… 141
10.6.1 广播 ……………………………………………………………………… 141
10.6.2 AllReduce ………………………………………………………………… 143
10.7 总结 ……………………………………………………………………………… 145
第 11章 使用 ROCm进行机器学习 …………………………………………
11.1 ROCm上的 PyTorch ……………………………………………………………… 146
11.1.1 安装 PyTorch …………………………………………………………… 147
11.1.2 测试 PyTorch安装 ……………………………………………………… 147
11.1.3 使用 Inception v3模型进行图像分类…………………………………… 148
11.2 ROCm上的 TensorFlow ………………………………………………………… 150
11.2.1 安装 TensorFlow ………………………………………………………… 150
11.2.2 测试 TensorFlow安装 …………………………………………………… 151
目 录
11.2.3 使用 TensorFlow进行训练 ……………………………………………… 151
11.3 总结 ……………………………………………………………………………… 152
第 12章 数据中心中的 ROCm…………………………………………………
12.1 容器化的 ROCm ………………………………………………………………… 153
12.2 使用 Kubernetes管理 ROCm容器 ……………………………………………… 154
12.3 使用 SLURM管理 ROCm节点 ………………………………………………… 156
12.3.1 SLURM交互模式 ……………………………………………………… 157
12.3.2 SLURM批处理提交模式 ……………………………………………… 157
12.4 总结 ……………………………………………………………………………… 158
第 13章 第三方工具 ……………………………………………………………
13.1 PAPI ……………………………………………………………………………… 159
13.1.1 PAPI工具和测试 ………………………………………………………… 160
13.1.2 PAPI对 AMD GPU的支持 ……………………………………………… 163
13.1.3 预设事件与计数器分析工具包 ………………………………………… 164
13.2 Score-P和 Vampir ………………………………………………………………… 165
13.2.1 使用 Score-P进行跟踪 ………………………………………………… 166
13.2.2 Score-P的使用 …………………………………………………………… 166
13.2.3 对 Quicksilver应用程序的性能分析 …………………………………… 167
13.3 Trace Compass和 Theia ………………………………………………………… 169
13.4 TAU ……………………………………………………………………………… 174
13.4.1 使用 TAU分析 HIP程序性能 ………………………………………… 174
13.4.2 使用 TAU跟踪 HIP程序 ……………………………………………… 175
13.4.3 使用 APEX测量 HIP程序 ……………………………………………… 178
13.4.4 TAU总结 ………………………………………………………………… 180
13.5 TotalView调试器 ………………………………………………………………… 180 13.6 HPCToolkit………………………………………………………………………… 184
13.6.1 HPCToolkit的工作流程 ………………………………………………… 185
13.6.2 使用 HPCToolkit分析 PIConGPU ……………………………………… 186
13.6.3 收集和分析性能数据和跟踪文件 ……………………………………… 186
13.6.4 使用硬件计数器进行测量 ……………………………………………… 188
目 录
13.7 使用 Linaro Forge进行调试和性能分析………………………………………… 190
13.7.1 Linaro DDT ……………………………………………………………… 190
13.7.2 Linaro MAP ……………………………………………………………… 191
13.7.3 Linaro性能报告 ………………………………………………………… 191
13.7.4 使用 Linaro DDT进行 GPU调试 ……………………………………… 192
13.7.5 使用 Linaro MAP分析 GPU性能 ……………………………………… 195
13.7.6 GPU性能报告 …………………………………………………………… 197 13.8 E4S——极端规模科学软件栈 …………………………………………………… 198
第 14章 HIP编程实战：大模型自注意力机制与 FlashAttention …………
14.1 自注意力机制简介 ……………………………………………………………… 200
14.2 FlashAttention核心思想 ………………………………………………………… 201
14.3 HIP实现 FlashAttention ………………………………………………………… 202
附录 A ROCm安装 ……………………………………………………………
A.1 先决条件…………………………………………………………………………… 216
A.2 理解 ROCm 软件包 ……………………………………………………………… 217
A.3 安装………………………………………………………………………………… 218
A.3.1 安装脚本方法 …………………………………………………………… 218
A.3.2 包管理器方法 …………………………………………………………… 219
A.3.3 验证安装过程 …………………………………………………………… 220
A.4 升级 ROCm………………………………………………………………………… 220
A.5 卸载 ROCm………………………………………………………………………… 221
附录 B CDNA汇编………………………………………………………………
B.1 使用 CDNA汇编代码 …………………………………………………………… 222
B.1.1 检索 HIP内核二进制文件 ……………………………………………… 222
B.1.2 反汇编 CDNA二进制文件 ……………………………………………… 223
B.2 CDNA寄存器 ……………………………………………………………………… 223
B.3 指令类型 …………………………………………………………………………… 224
B.4 内存访问指令 ……………………………………………………………………… 224
B.5 示例：移位复制 …………………………………………………………………… 225
B.6 示例：分支 ………………………………………………………………………… 226
目 录
B.7 比较 CDNA2 和 CDNA3 ………………………………………………………… 228
B.8 总结 ………………………………………………………………………………… 229
附录 C OmniTools ………………………………………………………………
C.1 Omnitrace ………………………………………………………………………… 230
C.1.1 Omnitrace 配置文件 ……………………………………………………… 230
C.1.2 收集跟踪数据 …………………………………………………………… 231
C.1.3 输出和可视化 …………………………………………………………… 232
C.2 Omniperf …………………………………………………………………………… 233
C.2.1 使用 Omniperf 分析程序性能 …………………………………………… 233
C.2.2 使用 CLI 进行分析 ……………………………………………………… 234
C.2.3 使用基于 Web 的 GUI 进行分析 ………………………………………… 236
C.2.4 使用 Grafana 进行分析 …………………………………………………… 236
C.3 总结 ………………………………………………………………………………… 238
参考文献……………………………………………………………………………
