HelloGPU-P00-入门与硬件速通

本系列学习笔记源自开源项目：https://github.com/datawhalechina/hello-gpu 。是自己在学习过程上的记录与心得。

第0篇入门与硬件速通

0.0 项目介绍与学习提示

阶段	你会学到	学完能做什么
入门（第 0 篇）	RX 9070 XT 环境、GPU 体系结构速通、第一个程序	跑通环境，建立 Roofline 心智模型
Profiling（第 1 篇）	benchmark、rocprof、性能记录	用数据说明算子慢在哪里
算子 + 刷题（第 2 篇）	Reduction / Softmax / GEMM / Flash Attention + 刷题方法论	把一个 naive kernel 一步步优化到接近硬件极限
Agent 算子层（第 3 篇）	Agent 入门、工具封装、多轮优化	让 Agent 生成、测量、裁决并记录候选，实现可复盘的优化循环
真实模型 + Agent（第 4 篇）	YOLO / LLM 部署 + Agent 自动优化	让 Agent 对真实模型做性能优化

上表是此次课程的所有学习内容。

0.1 环境准备

本次学习完成的事情主要是：准备学习所用的线上环境、验证学习环境中的软硬件环境是否可用。

0.1.1 使用AUP Learning Cloud环境

根据环境修改配置

此处使用“AUP Learning Cloud”的环境，账号申请通过后，登录启动：

启动学习环境后，只能以jovyan用户操作，进入如下目录/home/jovyan/hello-gpu/code/part0-intro，在左侧目录树中也打开此目录，双击pyproject.toml文件打开它进行编辑（因为无法以jovyan用户安装工具，所以无法安装vim之类的工具），将如下这些两处进行注释或修改。修改处1:（猜测：应该是容器有代理且能正常访问外网）国内镜像站（清华源/北外源）对这个相关翻墙代理/IP 返回了403 阻断/禁封，或者镜像站本身的防刷机制触发了封禁，故直接使用默认的 Python 官方源（PyPI）。将如下行注释掉。

...
# 大概第21行处
#[[tool.uv.index]]
#name = "pypi-mirror"
#url = "https://mirrors.bfsu.edu.cn/pypi/web/simple"
#default = true
...

修改处2:参考附录 B · 换一张卡：切换 ROCm 10.0 的 GPU 架构根据自己使用GPU型号与构架修改device-gfx1201 为合适的 “AMD GPU 硬件架构”的依赖扩展标记。此次使用的是，故修改成如下（即将第8行附近的device-gfx1201全部替换成device-gfx1151，共有3处需要替换）：

[project]
name = "part0-intro"
version = "0.1.0"
requires-python = ">=3.12,<3.13"
dependencies = [
    "numpy>=2.4.4",
    "matplotlib>=3.9",
    "torch[device-gfx1151]==2.13.0+rocm10.0.0",
    "torchvision[device-gfx1151]==0.28.0+rocm10.0.0",
    "torchaudio==2.11.0.2+rocm10.0.0",
    "rocm[devel,device-gfx1151]==10.0.0",
]

同步本篇 uv 环境：

1	`(part0-intro) jovyan@jupyter-sanyinjiang:~/hello-gpu/code/part0-intro$ uv sync`

同步完成后，激活本篇环境：

(part0-intro) jovyan@jupyter-sanyinjiang:~/hello-gpu/code/part0-intro$ source ./activate-rocm.sh
Activated ROCm uv environment:
  PROJECT:    /home/jovyan/hello-gpu/code/part0-intro
  VENV:       /home/jovyan/hello-gpu/code/part0-intro/.venv
  ROCM_PATH:  /home/jovyan/hello-gpu/code/part0-intro/.venv/lib/python3.12/site-packages/_rocm_sdk_devel
  Python:     /home/jovyan/hello-gpu/code/part0-intro/.venv/bin/python
  ROCm SDK:   10.0.0
Python 3.12.3

激活后检查 SDK 和 PyTorch 版本（gfx1151 / ROCm 10.0）：

(part0-intro) jovyan@jupyter-sanyinjiang:~/hello-gpu/code/part0-intro$ rocm-sdk version
10.0.0
(part0-intro) jovyan@jupyter-sanyinjiang:~/hello-gpu/code/part0-intro$ python -c "import torch; print(torch.__version__)"
2.13.0+rocm10.0.0

验证 GPU 可见性

(part0-intro) jovyan@jupyter-sanyinjiang:~/hello-gpu/code/part0-intro$ rocminfo | grep -E "^[[:space:]]*(Name|Marketing Name|Vendor Name|Device Type|Compute Unit):" | head -20
  Name:                    AMD RYZEN AI MAX+ PRO 395 w/ Radeon 8060S
  Marketing Name:          AMD RYZEN AI MAX+ PRO 395 w/ Radeon 8060S
  Vendor Name:             CPU                                
  Device Type:             CPU                                
  Compute Unit:            32                                 
  Name:                    gfx1151                            
  Marketing Name:          AMD Radeon 8060S Graphics          
  Vendor Name:             AMD                                
  Device Type:             GPU                                
  Compute Unit:            40                                 
      Name:                    amdgcn-amd-amdhsa--gfx1151         
      Name:                    amdgcn-amd-amdhsa--gfx11-generic

rocminfo 识别到 gfx1151 GPU（AMD Ryzen AI Max+ 395 (Radeon 8060S)）。

查看 GPU 的当前状态：

(part0-intro) jovyan@jupyter-sanyinjiang:~/hello-gpu/code/part0-intro$ amd-smi version
AMDSMI Tool: 27.0.0+6b0e43f3 | AMDSMI Library version: 27.0.0 | ROCm version: 10.0.0 | amdgpu version: N/A | ionic version: N/A
(part0-intro) jovyan@jupyter-sanyinjiang:~/hello-gpu/code/part0-intro$ amd-smi monitor
GPU  XCP  POWER   GPU_T   MEM_T   GFX_CLK   GFX%   MEM%   ENC%   DEC%      VRAM_USAGE
  0    0    7 W   24 °C     N/A   610 MHz    0 %    N/A    N/A    N/A    0.1/ 64.0 GB

验证 PyTorch ROCm

(part0-intro) jovyan@jupyter-sanyinjiang:~/hello-gpu/code/part0-intro$ python chapter1/check_torch_rocm.py
python: 3.12.3
torch: 2.13.0+rocm10.0.0
cuda_available: True
device_count: 1
device_name: AMD Radeon 8060S Graphics
result_shape: (1024, 1024)
result_dtype: torch.float32
result_device: cuda:0
result_checksum: 28381.671875

这个脚本做了三件事：导入 PyTorch、检查 GPU 后端、在 GPU 上跑一次 1024 × 1024 矩阵乘。此处出现的cuda相关字样，并不代表底层使用NVIDIA GPU，而是由于历史命名原因，将加速设备名称前缀写成固定的了，而不是根据不同的设备厂商来读取或映射不同的设备名称。

验证最小 HIP 程序

先确认一下 hipcc 版本

# 这个命令顺带还能验证编译器路径是否在 `_rocm_sdk_devel` 下面
(part0-intro) jovyan@jupyter-sanyinjiang:~/hello-gpu/code/part0-intro$ hipcc --version
HIP version: 7.15.26333-0000000
AMD clang version 23.0.0git (https://github.com/ROCm/llvm-project.git 8f497e0992fb7513f7f78a6f6b6f1056c375e961)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /home/jovyan/hello-gpu/code/part0-intro/.venv/lib/python3.12/site-packages/_rocm_sdk_devel/lib/llvm/bin

然后编译 HIP 程序：

# 检验 HIP 编译器能不能找到 GPU 对应的 device library
(part0-intro) jovyan@jupyter-sanyinjiang:~/hello-gpu/code/part0-intro$ cd chapter1/
(part0-intro) jovyan@jupyter-sanyinjiang:~/hello-gpu/code/part0-intro/chapter1$ hipcc vector_add.hip -O2 -o vector_add && echo "compile_status: PASS"
compile_status: PASS

运行程序：

(part0-intro) jovyan@jupyter-sanyinjiang:~/hello-gpu/code/part0-intro/chapter1$ ./vector_add
device_name: AMD Radeon 8060S Graphics
vector_size: 1048576
blocks: 4096
threads_per_block: 256
max_error: 0
status: PASS

0.1.2 使用AMD Radeon Cloud环境（优先推荐）

根据环境修改配置

本来要访问https://developer.amd.com.cn/radeon/创建容器实例。但此处很可能遇到资源不够而无法创建容器实例成功的现象，要根据https://zcnijjcepfie.feishu.cn/docx/F6vfdGhHroiNj2xPEJmcbKhBnzf的指导提交邮箱、使用自己创建的模板创建容器实例，创建成功后进入容器实例时如下图：

阅读后关闭上述chapter0.ipynb，打开Terminal，并进入目录/workspace/notebooks/part0-intro：

1
2
3

# bash
root@j-wznpfx43tvz7ipt7pccbiqifit:/workspace/notebooks/part0-intro/chapter0# cd ../
root@j-wznpfx43tvz7ipt7pccbiqifit:/workspace/notebooks/part0-intro#

容器中默认使用的命令行解释器是sh（全路径为：/usr/bin/sh），本人很不习惯。所以执行bash命令修改命令行解释器为bash（全路径为：/usr/bin/bash）。然后切换目录到/workspace/notebooks/part0-intro，执行后续操作。

另外，参考附录 B · 换一张卡：切换 ROCm 10.0 的 GPU 架构根据自己使用GPU型号与构架修改device-gfx1201 为合适的 “AMD GPU 硬件架构”的依赖扩展标记。此次使用的是“AMD Radeon Pro W7900D”，其构架是“device-gfx1100”，故修改成如下（即将第8行附近的device-gfx1201全部替换成device-gfx1151，共有3处需要替换）：

[project]
name = "hello-gpu-notebook-part0-intro"
version = "0.2.0"
requires-python = ">=3.12,<3.13"
dependencies = [
    #"torch==2.13.0+rocm10.0.0",
    "rocm[devel]==10.0.0",
    "hello-gpu-notebooks",
    "ipykernel>=7.4.0",
    "jupyterlab>=4.6.4",
    "nbclient>=0.11.0",
    "nbformat>=5.11.1",
    "pytest>=9.1.1",
    "triton>=3.8.0",
    "numpy>=2.5.3",
    "torch[device-gfx1100]==2.13.0+rocm10.0.0",
    "torchvision[device-gfx1100]==0.28.0+rocm10.0.0",
    "rocm[devel,device-gfx1100]==10.0.0",
]

同步本篇 uv 环境：

1	`root@j-wznpfx43tvz7ipt7pccbiqifit:/workspace/notebooks/part0-intro# uv sync`

这个过程耗时有点久，大概要10～20分钟。

同步完成后，激活本篇环境：

root@j-wznpfx43tvz7ipt7pccbiqifit:/workspace/notebooks/part0-intro# source ./activate-rocm.sh
Activated ROCm uv environment:
  PROJECT:    /workspace/notebooks/part0-intro
  VENV:       /workspace/notebooks/part0-intro/.venv
  ROCM_PATH:  /workspace/notebooks/part0-intro/.venv/lib/python3.12/site-packages/_rocm_sdk_devel
  Python:     /workspace/notebooks/part0-intro/.venv/bin/python
  ROCm SDK:   10.0.0
Python 3.12.3

激活后检查 SDK 和 PyTorch 版本（gfx1151 / ROCm 10.0）：

root@j-wznpfx43tvz7ipt7pccbiqifit:/workspace/notebooks/part0-intro# rocm-sdk version
10.0.0
root@j-wznpfx43tvz7ipt7pccbiqifit:/workspace/notebooks/part0-intro# python -c "import torch; print(torch.__version__)"
2.13.0+rocm10.0.0

验证 GPU 可见性

这是第一道门：ROCm 能看到 GPU。

root@j-wznpfx43tvz7ipt7pccbiqifit:/workspace/notebooks/part0-intro# rocminfo | grep -E "^[[:space:]]*(Name|Marketing Name|Vendor Name|Device Type|Compute Unit):|Version" | head -20
  Name:                    AMD EPYC 9334 32-Core Processor    
  Marketing Name:          AMD EPYC 9334 32-Core Processor    
  Vendor Name:             CPU                                
  Device Type:             CPU                                
  Compute Unit:            64                                 
  Name:                    AMD EPYC 9334 32-Core Processor    
  Marketing Name:          AMD EPYC 9334 32-Core Processor    
  Vendor Name:             CPU                                
  Device Type:             CPU                                
  Compute Unit:            64                                 
  Name:                    gfx1100                            
  Marketing Name:          AMD Radeon Pro W7900D              
  Vendor Name:             AMD                                
  Device Type:             GPU                                
  Compute Unit:            96                                 
      Name:                    amdgcn-amd-amdhsa--gfx1100         
      Name:                    amdgcn-amd-amdhsa--gfx11-generic

rocminfo 识别到 gfx1100 GPU（AMD Radeon Pro W7900D）。

查看 GPU 的当前状态：

root@j-wznpfx43tvz7ipt7pccbiqifit:/workspace/code/part0-intro# amd-smi version
AMDSMI Tool: 27.0.0+6b0e43f3 | AMDSMI Library version: 27.0.0 | ROCm version: 10.0.0 | amdgpu version: 6.16.13 | ionic version: N/A

root@j-wznpfx43tvz7ipt7pccbiqifit:/workspace/code/part0-intro# amd-smi monitor
GPU  XCP  POWER   GPU_T   MEM_T   GFX_CLK   GFX%   MEM%   ENC%   DEC%      VRAM_USAGE
  0    0   11 W   29 °C   24 °C    23 MHz    2 %    0 %    N/A    0 %    0.0/ 48.0 GB

验证 PyTorch ROCm

这是第二道门：PyTorch ROCm。

root@j-wznpfx43tvz7ipt7pccbiqifit:/workspace/code/part0-intro# python chapter1/check_torch_rocm.py
python: 3.12.3
torch: 2.13.0+rocm10.0.0
cuda_available: True
device_count: 1
device_name: AMD Radeon Pro W7900D
result_shape: (1024, 1024)
result_dtype: torch.float32
result_device: cuda:0
result_checksum: -22182.687500

验证最小 HIP 程序

这是第三道门：最小 HIP 程序。

先确认一下 hipcc 版本：

# 这个命令顺带还能验证编译器路径是否在 `_rocm_sdk_devel` 下面
root@j-wznpfx43tvz7ipt7pccbiqifit:/workspace/code/part0-intro# hipcc --version
HIP version: 7.15.26333-0000000
AMD clang version 23.0.0git (https://github.com/ROCm/llvm-project.git 8f497e0992fb7513f7f78a6f6b6f1056c375e961)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /usr/local/lib/python3.12/dist-packages/_rocm_sdk_devel/lib/llvm/bin

然后编译 HIP 程序：

# 检验 HIP 编译器能不能找到 GPU 对应的 device library
root@j-wznpfx43tvz7ipt7pccbiqifit:/workspace/code/part0-intro# cd chapter1/

root@j-wznpfx43tvz7ipt7pccbiqifit:/workspace/code/part0-intro/chapter1# hipcc vector_add.hip -O2 -o vector_add && echo "compile_status: PASS"
compile_status: PASS

运行程序：

root@j-wznpfx43tvz7ipt7pccbiqifit:/workspace/code/part0-intro/chapter1# ./vector_add
device_name: AMD Radeon Pro W7900D
vector_size: 1048576
blocks: 4096
threads_per_block: 256
max_error: 0
status: PASS

这是对当前环境的相关信息总结： ROCm 版本：10.0.0 Python 版本：Python 3.12.3 PyTorch 版本：2.13.0+rocm10.0.0 uv 环境所在路径： Activated ROCm uv environment: PROJECT: /workspace/code/part0-intro VENV: /workspace/code/part0-intro/.venv ROCM_PATH: /workspace/code/part0-intro/.venv/lib/python3.12/site-packages/_rocm_sdk_devel Python: /workspace/code/part0-intro/.venv/bin/python ROCm SDK: 10.0.0

当前 GPU 架构：
运行日期：2026-10-14 15:10:03 UTC

可以执行以下脚本文件输出（可命名为check.sh，放在/workspace/code/part0-intro目录下）：

#!/usr/bin/bash

source ./activate-rocm.sh
echo ""
echo ""

echo "ROCm 版本："`cat /opt/rocm/.info/version`
echo "Python 版本："`python -V`
echo "PyTorch 版本："`python -c "import torch; print(torch.__version__)"`
echo "uv 环境所在路径："
echo "Activated ROCm uv environment:"
echo "  PROJECT:    ${ROCM_PROJECT_DIR}"
echo "  VENV:       ${ROCM_VENV}"
echo "  ROCM_PATH:  ${ROCM_PATH}"
echo "  Python:     $(command -v python)"
echo "  ROCm SDK:   $(rocm-sdk version)"
echo ""
echo "当前 GPU 架构："``
echo "运行日期："`date "+%F %H:%m:%S"` UTC

0.2 GPU 体系结构（上）：编程模型与 wavefront 执行

这个小节，以最简单的情况讲述了AMD并行计算的实现机制，我的感觉就是将单线程顺序执行的重复任务以一种合适的机制映射到支持并行计算的硬件设备上，辅以并行计算软件层面的实现机制以并行的形式来执行单线程顺序执行的重复任务。

0.2.1 GPU计算的基本流程

不论是NVIDIA GPU还是AMD GPU，使用GPU做（并行）计算的大致流程都是：程序从 CPU （只能由CPU发起）上开始运行，准备输入，把数据放到 GPU 可以访问的位置，调用 GPU 完成计算。GPU 算完后，CPU 再取回结果。

GPU执行计算的具体内容放在一个函数内，这个函数叫核函数（Kernel）。CPU 提交这次 GPU 计算的动作，叫启动（Launch）。一次启动可以让许多线程（Thread）执行这个核函数，可以简单把每个线程看作是核函数的一次具体执行，GPU上的并行计算机制创建多个线程并行执行这个核函数，每个线程中执行核函数时会获取属于自己的数据进行计算。

在使用GPU做计算的过程中，CPU端一般叫做主机端（Host），GPU端一般叫做设备端（Device）。

线程也不是显卡上专属于它的一颗小处理器，它是一个软件术语。我们可以向GPU申请创建多个线程（针对特定的GPU，有一个线程上限），GPU 再利用有限的硬件资源安排执行。提交了多少份工作，和某一时刻有多少份工作正在执行，是两回事。

0.2.2 线程的组织与执行

线程块与网格

HIP 把一组线程称为线程块（Block）。启动 kernel 时，我们指定“每块多少个线程”和“一共多少块”。同一次启动的全部线程块，合在一起叫网格（Grid）。

为什么要创建线程块这个概念与机制： - 一方面，硬件可以把不同块安排到可用的资源上，不必等整个任务的所有线程都能同时执行。 - 另一方面，同一块内的线程可以在需要时交换数据、一起等待某个步骤完成。例如，把一批数合成一个和时，块内线程就可能要协作。

因此，Block 是组织一组线程工作的单位，Grid 是这次启动的全部工作集合。

线程块内的 wavefront

为了方便资源管理与计算，其实一个线程块内还可以继续把线程划分为不同的波前（Wavefront，简称 Wave），这样每个Wavefront包含多个不同的线程，此时Wavefront中每个线程叫做一个通道（Lane）。一个 Block 里的多个 Wavefront 不必同时开始或结束。

为什么要创建wavefront这个概念与机制： - GPU 可以让同一 Wavefront 中当前参与执行的线程，共同执行一条指令，各自使用自己的数据。

分块和组成 wavefront 也有不同的作用：Block 规定哪些线程归在一起、可以协作；Wavefront 则把块内线程组织成共同执行指令的小组。

启动 kernel 时指定块数和每块线程数；GPU 按这个 kernel 采用的 wave 大小（这是个相对固定的值）组织执行，不需要再单独指定“启动几个 wavefront”。不同 Block 的线程不会拼成同一个 wavefront。

随堂问题：同一个 kernel 启动出许多线程，它们是各执行工作单的一步，还是各执行一遍工作单？一个 Block 有多个 wavefront，是否说明多出了新线程？回答：各执行一遍工作单，即执行核函数kernel中设定的逻辑；没有多出新线程，只是把几个线程划分为同一个Wavefront，方便它们执行相同的指令。

AMD中这套软件术语在NVIDIA中都有对应的术语，可以对比记忆：

HIP / CUDA 常见名称	AMD 语境中的称呼
thread	线程；底层也称 work-item
block / thread block	工作组 workgroup
warp	波前wavefront；大小按当前目标与 kernel 确认
grid	一次启动中的全部工作组
shared memory / `__shared__`	工作组共享的 LDS

要在AMD执行并行计算，需要以HIP语法定义核函数kernel（很像CUDA C编写的kernel），如下：

// GPU 内核函数：每个启动的 GPU 线程最多计算向量中的一个元素。
// `__global__` 表示此函数由 CPU 端（host）发起调用、在 GPU 端（device）并行执行。
// a、b 是输入向量，c 用来保存逐元素相加的结果，n 是向量的有效元素个数。
__global__ void vector_add(const float* a, const float* b, float* c, int n) {
  // 将三维线程索引在本例中使用的 x 维展开为全局一维元素下标：
  // 每个 block 有 `blockDim.x` 个线程，`blockIdx.x` 指明当前是第几个 block，
  // `threadIdx.x` 则是当前线程在本 block 内的编号。
  int idx = blockIdx.x * blockDim.x + threadIdx.x;

  // 最后一个 block 往往不能被 n 整除，部分线程的 idx 会越过向量末尾。
  // 此边界判断可防止这些多出的线程进行非法内存访问。
  if (idx < n) {
    // 每个合法线程只写入自己对应的 c[idx]，线程之间不会写同一位置，
    // 因此这里不需要额外的同步或原子操作。
    c[idx] = a[idx] + b[idx];
  }
}

核函数kernel只能由CPU发起调用，并在调用传入工作组/线程块个数，与每个工作组/线程块中线程数，如下：

1
2
3

// 使用 HIP/CUDA 风格的 <<<grid, block>>> 语法启动内核：`blocks` 是 grid 中的
// block 数，`threads` 是每个 block 的线程数。该启动通常是异步的，CPU 会继续执行。
vector_add<<<blocks, threads>>>(d_a, d_b, d_c, n);

0.2.3 线程编号与数据索引

0.2.3.1 线程块与块内线程的编号

既然每个线程在执行核函数时要读写属于自己的数据，那么就需要对每个线程进行区分。跟国家行政单位或军队编制单位等拥有明显的上下级组织机构类似，要先定位一个上级机构组织，才能定位到其下属的下级机构组织。

所以说，在AMD GPU编程中，要标识一个线程，要先标识一个线程块。对于一次启动创建的所有线程块与线程而言，每个线程块拥有一个唯一的编号，线程块内的每个线程拥有一个唯一的编号，但不同线程块的线程编号从0开始、所以可能会相同。

如果给定一个线程块编号，并给定此线程块编号对应线程块中的线程编号，就可以唯一确定一次启动中创建的所有线程中某一个线程，这样就可以把将所有线程区分开来了。

以下面这个示例来讲述相关机制：现有两个长度为10的整数数组a、b，需要将这两个数据中对应索引位置的两个元素两两相加，并将结果保存到另一个数组c中。因为这种两两相加互不依赖与影响，完全可以用AMD GPU来计算。

如上图所示，三个线程块中都有一个 0 号线程，所以只说“0 号线程”还不够。说成“第 1 个线程块中的 0 号线程”，才能指出具体是哪一个。

启动时，CPU 端代码指定多少块、每块多少线程；HIP 让每个线程读取自己的块号和块内编号（有特定的实现机制），不需要我们逐个给线程填写编号。

0.2.3.2 数组下标的计算

本例选择一种简单的分工：每块负责连续的 4 个数组位置，块内的 0～3 号线程分别处理这四个位置。 第 0 块负责下标 0～3，第 1 块负责下标 4～7，后面的块依次接上。我们先跟着第 1 块里的 2 号线程，看它怎样找到自己的位置。

第 1 块负责的这一段从下标 4 开始，块内的 2 号线程再偏移两个位置，因此负责下标 4 + 2 = 6。这里的起点 4，又可以写成“块号 1 × 每块 4 个线程”。

可以进行归纳，把同一条分工规则写给所有线程，就是： 本例的数组下标 = 块号 × 每块线程数 + 块内线程编号。

HIP 为刚才用到的三个量提供了对应的名字（就是我上面所说的：“HIP 让每个线程读取自己的块号和块内编号（有特定的实现机制）”）：

代码里的名字	含义	选中线程读到的值
`blockIdx.x`	我属于第几块	1
`blockDim.x`	每块有多少个线程	4
`threadIdx.x`	我在本块里是第几个线程	2

这样一来，如果要用（伪）编码形式给每个线程一个唯一编号来区分或唯一确定一个线程，就是如下：

1 2	`# 线程唯一编号 = 块号 × 每块线程数 + 块内线程编号 int idx = blockIdx.x * blockDim.x + threadIdx.x;`

每个线程取得这样一个唯一编号后，执行kernel同一套逻辑：c[idx] = a[idx] + b[idx] ，就可以把“并行把两个数组a、b相同索引对应元素两两相加、结果保存到另一个数组c对应索引位置上”这个功能实现了。

每个线程能确定自己要读写的数据即可，并不一定就是idx=0的线程一定只能处理数组a、b中索引为0的元素，也不一定是idx=0的线程将计算结果保存到数组c中索引为0的位置上。其中的映射逻辑是灵活的，但需要充分利用好所有创建的线程，分别把所有要处理的数据覆盖并且规避非法或异常情况，即保证分工没有遗漏、重复写入或越界。

线程块号及块内线程号，并不表示执行顺序。即第 1 块不需要等第 0 块算完，才能计算自己的；第 1 块内线程 3 也不需要等线程 0 算完才算自己的。

0.2.3.3 数组边界检查

如前所述，整数数组a、b的长度假设都为10，有效下标是 0～9。每块仍安排 4 个线程，两个块只能覆盖 8 个位置，因此需要第三块。

本来三个块应该是对应12个线程，但数组a、b、c都只有10个元素，按照我们的设计，每一个线程读写a、b、c中唯一个索引对应的元素，这样一来第三块中最后两个线程就是多余的了，应该让它们什么都不干。那首先要把它们找出来，这个找出来的过程中就涉及到数组边界检查。

如图所示，块2内的 0、1 号线程算出下标 8、9，正好处理最后两个元素。块内的 2、3 号线程算出 10、11，已经超出数组范围。我们可以用 if (idx < n) 检查下标：只有条件成立时，才执行数组读写。（此处 n 是数组a、b的长度。且边界检查的情况可能比较复杂，这里可以通过简单的条件判断达到目的。）

把“算下标、查边界、做加法”全部一起实现，就是第 1 章的 Vector Add 源码实现的核心内容：

__global__ void vector_add(const float* a, const float* b, float* c, int n) {
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx < n) {
    c[idx] = a[idx] + b[idx];
  }
}

其中，__global__ 标记这是由 CPU 发起、在 GPU 上运行的 kernel；a、b、c 指向 GPU 上的数组，n 是要处理的元素个数。

读函数体时，始终站在一个线程的角度：我先算出自己负责的位置，检查没有越界，再读取两个数，把和写回。其他线程执行的是同一个函数，只是代入的编号和读取的数据不同。

0.2.3.4 线程编号与lane编号

下面两张图，是在演示这样一个示例：某次启动一共创建了两个线程块、每个线程块中包含64个线程。且对于当前所使用的GPU硬件，此处使用的是wave32的配置，就是说每个线程块会分成多个wavefront，规定每个wavefront包含32个线程（即32个通道）。这里每个wavefront包含32个线程上图中显示，块0内线程 0～31 属于第一个 wavefront，32～63 属于第二个；各小组（wavefront）里的位置编号（通道编号）都从 0 开始。

wavefront中每个线程都称为通道(lane) 此图告诉我们，wavefront中用通道(lane)来标识每个线程的位置。因为此处使用wave32配置，所以每个wavefront中包含32个通道，这32个通道的编号为 0~31 。

如果想要选中块0内线程 35：第二个小组从线程 32 开始，所以它的组内位置是 35 − 32 = 3，也就是 lane 3。块0内，“线程 35” 和 “wavefront1、lane 3” 描述的是同一份工作，以及它在小组里的位置。

0.2.3.5 线程分组与硬件执行单元

到目前为止，我们已经学习到 AMD GPU编号内线程相关的组织形式：“Grid → Block → Wavefront → Thread”，也知道了具体每个线程如何获取唯一标识读写属于自己的数据了。

一个 Block 或 wavefront 是一组线程工作，是软件层面的概念；要把它们与硬件设备区分开（虽然它们有联系），并不表示显卡上为一个 Block 或 wavefront 单独准备了一块同名硬件。

在 AMD RDNA 的语境里，先区分下面几个名字：

名字	本章先掌握的含义
单指令多数据执行单元（Single Instruction, Multiple Data，SIMD）	用一条向量指令处理一组 lane 的数据；这里的“多数据”对应各 lane 不同的输入
计算单元（Compute Unit，CU）	包含执行单元及相关资源的一层硬件组织
工作组处理器（Workgroup Processor，WGP）	RDNA 中包含 CU 的更外层组织，也与同一工作组的放置范围有关

下图展示了工作分组与硬件单元的关系。图没有规定单元数量或某次运行的调度顺序，也不表示一个 block 独占一个硬件单元。从上图可以看出（实际上也确实是这样）：软件上的一组工作需要落到硬件资源上执行。AMDGPU 文档 LLVM AMDGPU 执行模型对工作组给出了放置约束：同一工作组的 wavefront 在同一 WGP 中执行；更细的执行范围还与 CU/WGP 模式有关。详情见文档。

在进行接下来更细致的学习之前，先有这样的认知：逻辑上创建了多少线程，与物理实际上此刻能执行多少工作，是两个问题。 后者会受到其他很多资源（如寄存器、共享存储和数据等待等等）的制约。就是说，对于一份工作而言，当设定好了主机端的逻辑与设备端核函数逻辑后，某时刻GPU上真正能并行执行的线程数有一个理论上限，实际线程数会在这个理论线程数左右，无法做到有远超此理论线程数的线程在执行。

同一个工作组内的不同线程/通道也可以按条件决定是否执行。同一 wavefront 内的线程选择了不同控制路径，这种情况专业术语叫分支发散（Branch Divergence）。

0.2.4 本章总结

CPU 端程序准备数据、启动 kernel，并在使用结果前确认 GPU 已经完成计算。
Kernel 是共同的计算规则，Thread 是这份规则的一次执行。本例让每个有效线程处理一对数，每个线程都完成取数、相加和写回。
启动时指定块数和每块线程数，全部块组成 Grid；块内线程按 wave 大小组成执行小组。Lane 表示线程在 wavefront 内的位置。
理解分工后，编号才派上用场：HIP 提供线程身份，kernel 中的索引规则决定数据位置，边界检查保证读写有效。
Block 和 wavefront 是线程的分组，CU、WGP 和 SIMD 是承接这些工作的硬件。提交的线程总数不等于某一时刻正在执行的线程数。

0.2.5 自我检验

用“每块 4 个线程、共 10 个元素”的手算配置，指出 block 1 的线程 3 读写哪个位置。
在一个 256 线程的 block、wave32 的配置下，指出块内线程 63 属于哪个 wavefront、对应哪个 lane。
wavefront 0 全走 A、wavefront 1 全走 B，是否已经构成了 wavefront 内发散？说明判断依据。
CPU 启动一次数组加法 kernel，里面有许多线程。Kernel 和 Thread 分别指什么？每个线程都要取数、相加、写回吗？
用“营 → 连 → 排 → 单兵”说出 Grid、Block、Wavefront、Thread 的包含关系。Lane 表示什么？哪些分组数量由启动代码指定？
按本例连续分工、每块 4 个线程处理 10 个元素，第 1 块的线程 3 负责哪里？第 2 块的线程 3 为什么要跳过？

回答：

blockIdx.x为1，blockDim.x为4，threadIdx.x为3。可知idx = blockIdx.x * blockDim.x + threadIdx.x = 1 * 4 + 3 = 7。
wave32表示每个wavefront中有32个线程，又知每个线程组中有256个线程。可知：一共有256/32=8个wavefront，在当前工作组内，线程0~31属于wavefront1，线程32～63属于wavefront2，它对应的lane编号是63取余32即63%32=31，即lane31 。
没有构成。同一 wavefront 内的线程选择了不同控制路径才叫wavefront 内发散，不同的wavefront无法比较与确定。
Thread 是线程，一次启动会创建多个线程块，每个线程块内包含多个线程。Kernel 是核函数，表示GPU上每个线程的计算规则。多个线程的计算规则是一样的，只是读写的数据不同，分别完成不同的任务，等到所有线程完成，整个并行计算任务才算完成。在数组相加示例中，每个线程都要取数、相加、写回。
营 → 连 → 排 → 单兵，分别代表 Grid、Block、Wavefront、Thread，它们一级嵌套一级的上下级与包含关系，一个营（Grid）包含多个连（Block），一个连（Block）包含多个排（Wavefront），一个排（Wavefront）包含多个单兵（Thread）。Lane 表示wavefront 中线程的位置。线程块个数、每个线程块内包含的线程数由启动代码指定，每个wavefront 内包含的通道数由硬件本身决定，一般来说是稳定的。
第 1 块的线程 3 负责的数组元素索引是：1（表示块编号） * 4（每个线程块的线程数） + 3（wavefront内线程位置号，即lane） = 7 。第 2 块的线程 1 处理的数组元素索引是 9，数组元素索引最大值是9，所以第 2 块的线程 1 处理就是数组中最后一对元素，那么第 2 块的线程 1 后面的线程全部不需要干活，所以要跳过。如果它们要强制去读写数据，会造成数组下标越界，非法、产生不可预料的后果。

0.2.6 完整做vector_add的HIP示例

最后给出一个完整可执行的HIP示例：

// HIP Runtime API：负责 GPU 内存分配、主机/设备间数据传输、内核启动后的错误查询等。
// 在 AMD ROCm 平台上，HIP 提供了与 CUDA 相近的编程模型和接口风格。
#include <hip/hip_runtime.h>

// `std::abs` 和 `std::max` 分别用于计算数值误差的绝对值和最大值。
#include <cmath>
// 使用标准输出流打印设备信息和验证结果。
#include <iostream>
// 使用 `std::vector` 管理主机（CPU）端数组，自动负责其内存生命周期。
#include <vector>

// 将每个 HIP Runtime API 调用包裹为统一的错误检查。
//
// `call` 会先被执行，其返回的 `hipError_t` 状态码保存到 `err`。若状态不是
// `hipSuccess`，就打印可读的错误信息并从 `main` 返回失败码。`do { ... } while (0)`
// 使这个多行宏在语法上表现得像一条普通语句，因此可安全地写成 `HIP_CHECK(...);`。
#define HIP_CHECK(call)                                                    \
  do {                                                                     \
    hipError_t err = call;                                                 \
    if (err != hipSuccess) {                                               \
      std::cerr << "HIP error: " << hipGetErrorString(err) << std::endl;  \
      return 1;                                                            \
    }                                                                      \
  } while (0)

// GPU 内核函数：每个启动的 GPU 线程最多计算向量中的一个元素。
// `__global__` 表示此函数由 CPU 端（host）发起调用、在 GPU 端（device）并行执行。
// a、b 是输入向量，c 用来保存逐元素相加的结果，n 是向量的有效元素个数。
__global__ void vector_add(const float* a, const float* b, float* c, int n) {
  // 将三维线程索引在本例中使用的 x 维展开为全局一维元素下标：
  // 每个 block 有 `blockDim.x` 个线程，`blockIdx.x` 指明当前是第几个 block，
  // `threadIdx.x` 则是当前线程在本 block 内的编号。
  int idx = blockIdx.x * blockDim.x + threadIdx.x;

  // 最后一个 block 往往不能被 n 整除，部分线程的 idx 会越过向量末尾。
  // 此边界判断可防止这些多出的线程进行非法内存访问。
  if (idx < n) {
    // 每个合法线程只写入自己对应的 c[idx]，线程之间不会写同一位置，
    // 因此这里不需要额外的同步或原子操作。
    c[idx] = a[idx] + b[idx];
  }
}

int main() {
  // 向量长度为 2^20（1,048,576）个单精度浮点元素。
  // 使用 2 的幂便于演示大规模数据上的并行计算，但不是算法的必需条件。
  const int n = 1 << 20;
  // 一个 float 占 `sizeof(float)` 字节；该值用于设备端分配和数据拷贝。
  const size_t bytes = n * sizeof(float);

  // 在主机内存中创建输入、输出向量。构造函数的第二个参数会初始化所有元素：
  // a 全为 1、b 全为 2，故正确的逐元素相加结果应当全部为 3。
  std::vector<float> h_a(n, 1.0f);
  std::vector<float> h_b(n, 2.0f);
  std::vector<float> h_c(n, 0.0f);

  // d_ 前缀表示设备（device）端指针。先初始化为 nullptr，便于明确其尚未分配 GPU 内存。
  float* d_a = nullptr;
  float* d_b = nullptr;
  float* d_c = nullptr;

  // 在 GPU 全局内存中为两个输入向量和一个输出向量各分配 `bytes` 字节空间。
  // 每次调用均经 HIP_CHECK 检查，例如显存不足时会立刻报错并退出。
  HIP_CHECK(hipMalloc(&d_a, bytes));
  HIP_CHECK(hipMalloc(&d_b, bytes));
  HIP_CHECK(hipMalloc(&d_c, bytes));

  // 将 CPU 端准备好的两个输入向量复制到对应的 GPU 内存。
  // `hipMemcpyHostToDevice` 明确指示传输方向，避免 HIP 猜测指针所属位置。
  HIP_CHECK(hipMemcpy(d_a, h_a.data(), bytes, hipMemcpyHostToDevice));
  HIP_CHECK(hipMemcpy(d_b, h_b.data(), bytes, hipMemcpyHostToDevice));

  // 每个 block 启动 256 个线程。block 数采用向上取整，确保至少为 n 个元素
  // 启动足够多的线程；最后一个未填满的 block 由内核中的边界判断处理。
  const int threads = 256;
  const int blocks = (n + threads - 1) / threads; // 4096

  // 使用 HIP/CUDA 风格的 <<<grid, block>>> 语法启动内核：`blocks` 是 grid 中的
  // block 数，`threads` 是每个 block 的线程数。该启动通常是异步的，CPU 会继续执行。
  vector_add<<<blocks, threads>>>(d_a, d_b, d_c, n);

  // `hipGetLastError` 检查内核“启动”本身是否失败（例如配置参数不合法）。
  // `hipDeviceSynchronize` 等待 GPU 完成实际计算，并报告执行期间出现的异步错误。
  HIP_CHECK(hipGetLastError());
  HIP_CHECK(hipDeviceSynchronize());

  // GPU 计算结束后，将结果向量从设备端复制回 CPU 端，以便进行数值验证。
  HIP_CHECK(hipMemcpy(h_c.data(), d_c, bytes, hipMemcpyDeviceToHost));

  // 遍历结果，计算所有元素与理论值 3.0f 之差的最大绝对值。
  // 这里输入值 1、2 和结果 3 都可被 float 精确表示，因而预期最大误差正好为 0。
  float max_error = 0.0f;
  for (int i = 0; i < n; ++i) {
    max_error = std::max(max_error, std::abs(h_c[i] - 3.0f));
  }

  // 查询编号为 0 的设备属性，用于在输出中标识实际执行计算的 GPU。
  // 代码成功完成前已运行过 GPU 内核，因此此处可确认至少存在一个可用设备。
  hipDeviceProp_t prop{};
  HIP_CHECK(hipGetDeviceProperties(&prop, 0));

  // 打印实验规模、启动配置和验证结论，便于快速确认运行环境及计算是否正确。
  std::cout << "device_name: " << prop.name << std::endl;
  std::cout << "vector_size: " << n << std::endl;
  std::cout << "blocks: " << blocks << std::endl;
  std::cout << "threads_per_block: " << threads << std::endl;
  std::cout << "max_error: " << max_error << std::endl;
  std::cout << "status: " << (max_error == 0.0f ? "PASS" : "FAIL") << std::endl;

  // 释放前面由 hipMalloc 获取的三块设备内存，避免长时间运行的程序造成显存泄漏。
  // `std::vector` 管理的主机内存则会在 `main` 结束时自动释放。
  HIP_CHECK(hipFree(d_a));
  HIP_CHECK(hipFree(d_b));
  HIP_CHECK(hipFree(d_c));

  // 将数值验证结论映射为进程退出码：0 表示成功，1 表示检测到计算结果异常。
  return max_error == 0.0f ? 0 : 1;
}

0.3 GPU 体系结构（下）：片上资源与数据通路

本章从更细化、深入的角度来分析上一章节中两个数组相加的示例、学习核函数中数据存储与访问的具体过程；介绍了同一wavefront内不同线程进行合并访存的必要性与可行性；学习了同一个线程块内不同wavefront间线程利用LDS、屏障同步机制完成数据交接、进度对齐；最后学习不同wavefront 的执行情况可以被GPU调度器根据情况实时变更，实现执行资源利用率最大化。

0.3.1 数据位置与读写过程

还是以”两个数组a、b对应位置元素相加，结果保存到数组c中对应位置“为例：此图（图0.3.1-1）是两个数组相加示例的线程执行过程的粗略总体示意图。

两个数组a、b中都有多个元素（假设为n个）需要对应相加，理论上需要n个线程来完成。这里取其中下标为2的数组元素为例，细致分析这个线程中具体做的事情，如下图（图0.3.1-2）所示：

上述两个图中几个关键部位都有一个对应的专门术语：

全局内存（Global Memory）：图0.3.1-1中，存放输入与输出数组的位置（存储空间）。全局内存是从软件层面来讲的，全局内存中的数据实际由”GPU显卡“硬件来承载。
寄存器（Register）：图0.3.1-1中，存放临时值对应的位置。是指指令要使用到的临时存储空间，保存的是指令使用到的临时值（包括指令需要的输入与指令执行后的输出）。
执行单元：图0.3.1-1中，执行加法对应的逻辑与物理单元。它读入寄存器中的数据、执行指令、再将指令执行结果写入寄存器。

上述就是GPU内，执行c[2] = a[2] + b[2]那个线程做的事情。

0.3.2 Wavefront 的访存组织

前面说明过：同一个wavefront内不同lane对应的线程，执行完全相同的指令。就是说，对于同一个wavefront，当它执行一条数组读取指令时，这个wavefront内的所有线程分别有自己要读取与使用的数组元素，此时AMD HIP编程的实现是，去把这些线程要读取的数据所在数据块（每个数据块可能包含多个数组元素，根据自定义数据结构与数据类型而异）读取来，但实际上最后可能只使用其中一部分数据。

打个比方：某次任务中，其中一个wavefront包含多个线程，下图只画出其中4个。由于自定义数据结构与数据类型的不同，可能会出现下图中”相邻位置“、”分散位置“两种情况（可以把下图中的教学组看成一个数据块）：在GPU显存中，数据是按数据块进行组织与访问。一个数据块的大小一般来说是相对稳定的，就相当于格式化一块磁盘会将选定其数据块大小一样，所以说每个数据块中能够保存的数据是有限的。

即存储系统通常按一定大小的数据块组织和服务访问。线程需要一个元素，不意味着底层只处理这一个元素。请求集中时，同一段数据（一个或少数几个相邻数据块）就可能满足多个线程；请求分散时，可能需要涉及更多数据块，其中只有部分元素被这批线程使用。

对于”相邻位置“这种情况，因为自定义数据结构设计地科学合理（与单个数据块刚好匹配：可能是一样大，也可能是几个结构体只好凑成一个数据块等情况），同一个wavefront的多个线程需要的数据，存储在一个或少数几个数据块内，访问这一个或少数几个数据块即可。对于”分散位置“这种情况，同一个wavefront的多个线程需要的数据会更分散地存储在不同的数据块中（典型的情况是：一个数据块中只有很少一部分数据是线程所需要的，把整个数据块读取过来浪费了寄存器、带宽等资源），此时就需要读取多个数据块，造成临时存储空间与带宽等资源的浪费。

硬件把同一 wavefront 的访存请求尽量合并处理，这个过程称为合并访存（Memory Coalescing）。所以在此背景下，在GPU编程中，不仅仅是要“写 Kernel 函数”，更核心的功夫在于设计数据在显存中的排布方式（尽量连续与对齐），以及合理分配线程的工作任务（让相邻线程恰好去读连续的数据），从而把分散的访存请求压缩成尽量少的数据块传输，提高合并访存的效率。

访存地址由我们编写的程序决定，实际怎样合并由硬件完成。

数据读取还存在另一种机会：需要的数据可能已经有一份副本留在更靠近计算部件的存储中。这种由硬件管理的存储叫缓存（Cache）。找到所需副本叫命中（Hit），没找到则继续向后面的存储层级请求。缓存通常按数据块保存副本，这些数据块称为缓存行（Cache Line）。

合并访存关注的是一批请求所需数据的位置是否相邻、分布在一个或少数几个相邻的数据块中；缓存关注的是请求所需数据是否已在更靠近计算部件的存储中存在副本。

0.3.3 Block 内的数据共享与同步

上面的示例，各个线程只需要关注自己所需数据，且不需要互相协作或共享数据，实现上很简单。

现在给出一个数据简单但需要引入数据共享与同步机制的新示例：把 [1, 2, 10, 20] 中的四个数求成一个总和。

因为”在单个 wavefront 内部：硬件天然锁步，无需显式做屏障同步。“这样的机制存在，所以单个 wavefront 内部不存在屏障同步。触发屏障同步的线程理论上应该属于不同wavefront（如果线程块里只有 1 个 wavefront，此时调用 __syncthreads() 在逻辑上虽然能通过编译，但对于单个 wavefront 内部而言并没有实际的硬件等待意义）。

所以此示例假设需要用两个线程甲、乙来实现，且这两个线程分别属于同一个线程块的不同wavefront：wave0、wave1。

其中线程甲、线程乙分别计算 1与2之和、10与20之和，当两种线程得出自己结果（先后顺序不一定）后，分别将中间结果保存共享区LDS中。之后这两个线程在共同同步点会合、之后不同线程才能执行同步点之后的操作，并保证同步点前的共享写入对组内线程可见。

上图中，同一个线程块内线程甲需要读取线程乙保存的计算结果，线程乙内的普通局部变量不是甲可以直接读取的共享位置。为此需要在两个线程之外提供一个双方都能使用的交换区域。可以使用 AMD GPU 的局部数据共享存储（Local Data Share，LDS） 提供这个交换区域，在HIP中，用 __shared__ 声明供同一 block 的线程使用的共享存储。

上图中，线程甲把计算结果写入LDS中后，并不能保证线程乙也已经完成（实际情况是：线程甲先完成或后完成都可能发生）。此时需要引入一个同步机制，保证线程甲、线程乙（如果还存在其他线程，就需要等待并保证所有线程完成）都完成。这个同步机制就是屏障同步（Barrier）。 屏障同步（Barrier） 是一种多线程协调机制：它强制规定一组并发执行的实体（如不同的 wavefront）在到达该指定代码位置时必须暂停等待，直到所有参与者全部到达后，方可被共同放行继续向下执行，以此实现进度对齐与内存可见性/安全性。

HIP 的 __syncthreads() 让 block 中的线程在共同同步点会合，并保证屏障前的共享写入对组内线程可见。本例把屏障同步放在局部和写入之后、合并读取之前；block 中的全部线程都要到达它，不能让部分线程绕过。HIP 同步函数

等并发执行的线程甲、线程乙都到过屏障同步点后，线程甲再从共享区LDS中取出所有线程各自写入的数据，执行后续的求和操作。在这一过程里，LDS 提供交接位置，屏障确定何时可以使用交接的数据。

同样都是直接从中取值，缓存是由硬件管理，保存访问过的数据副本；而共享区LDS则是由程序在软件层面来安排，用于数据交接与共享。

0.3.4 访存等待与指令调度

一个wavefront被调度执行后，可能并不会一直占用系统的计算资源执行下去直到结束。因为它可能会在某个步骤发生读写请求，此时它申请读取或写入动作可能并不会立刻成功结束（比如申请读取数据时，但数据并不会立刻返回来），此时，如果后续步骤依赖此步骤的读写结束，wavefront就要等待，具体为是状态变更为“驻留”。

由此引发对wavefront的几种状态的探讨：

状态	可以怎样理解
驻留（Resident）	工作已经获得相应片上资源，但因为数据等未准备就绪无法被立刻调度以执行。但其状态保留在硬件中
就绪（Ready / Eligible）	满足继续执行所需的条件，可以被选择
执行（Executing）	当前获得了执行机会，正在发射或执行相应指令

回到上述4个数求总和的示例，wave0的线程甲读取数据过程，wave0 已经发出请求，但输入还没有返回，依赖它的加法暂时不能执行。如果同一计算资源上还有 wave1 的指令已经就绪，硬件可以先安排 wave1。

这种选择就绪指令放到执行资源上真正执行的过程由 GPU 的调度器（Scheduler） 发起。

wave0 等待时，它的临时值和执行状态仍然保留；选择 wave1 不需要把 wave0 当作已经完成，也不需要由 CPU 手动切换这两个 wavefront。

这种用其他就绪工作利用等待时间的方式，称为延迟隐藏（Latency Hiding）。wave0 的读取并没有因此更早返回，只是执行资源在这段时间有机会继续做别的工作。如果其他 wavefront 也都在等待，或者它们需要的执行资源正忙，仍然可能出现资源空闲。AMD GPUOpen：Occupancy explained

本章小结

输入、输出数组与计算使用的临时值各有用途。寄存器保存值，执行单元使用这些值运算；得到结果后仍要写回输出。
看访存要观察同一 wavefront 的地址排列。集中请求更有机会合并服务，缓存则可能提供已有的数据副本。
当后续计算需要其他线程的中间结果时，可以用 LDS 安排交接，并用正确的 block 同步保证先写后读。
一个 wavefront 等待数据时，调度器可以选择其他就绪指令；这利用了等待时间，但没有消除访问本身的延迟。

自我检验

加法已经得到临时结果 33，为什么 c[2] 还可能没有更新？执行加法与保存临时值分别由什么部件承担？
两组线程都只读取四个数，为什么不能仅凭读取数量判断访问工作量？
求和例子中的 30 从哪里来，由谁写入共享区，最后由谁读取？甲写完自己的位置后，为什么还需要同步？
调度器选择 B（wave1）执行时，等待中的 A（wave0）是否已经完成？A（wave0）的输入会因此更早返回吗？

回答：

临时结果是在执行单元上执行得到的，需要先保存到寄存器中，然后再更新到保存在全局内存中的数组具体位置上。执行单元、寄存器。
要看所需的四个数的地址分布，即它们是否在相同或相邻数据块上。如果在相同或相邻数据块上，则只读取一个或少量数据块即可，数据访问工作量少；如果在多个不同的数据块上，则需要读取多个不同的数据块，每个数据块中只取很少数据，数据访问工作量大。 另外，是否已有缓存也会影响访存工作量。
30 由wave1的线程乙方计算得到，并写入到LDS中，所有相关线程到达屏障同步后由wave0中的线程甲读取。线程甲写完自己的位置后，并不能保证其他线程已完成，这是不确定的，所以需要利用屏障产同步机制实现等待。
等待中的 A（wave0）未完成，它因等待的输入未到达而让出执行资源，保存了临时值与执行状态，并未完成。A（wave0）的输入不会因此更早返回。

0.3 GPU 体系结构（下）：片上资源与数据通路（旧）

0.3.1 一次加法的大致过程

AMD GPU上要执行一次计算，首先要从全局内存（显存。同一事物从不同角度的不同叫法）或缓存（如果能在缓存中找到对应的数据）读取数据加载到临时存储（一般是寄存器），然后把临时存储中的输入数据进行计算得到结果，再将结果写回到全局内存（显存）中。

另外在全局内存（显存）的外面，还有一个CPU使用的内存。就是说计算前有时需要从CPU使用的内存中将复制到全局内存（显存）中，计算完成后可能需要将全局内存（显存）中的结果再又复制回CPU使用的内存中。

0.3.2 数据的存取、管理与缓存

0.3.2.1 向量通用寄存器与标量通用寄存器

同一个 wavefront 下的不同lane可以存取不同的数据，就是说同一个 wavefront 下的不同lane 有自己独有存取数据的存储空间，这是通过 向量通用寄存器（Vector General-Purpose Register，VGPR） 来实现的。同一个 wavefront 下的不同lane 在同一个向量通用寄存器编号下有自己的独有的寄存器分量（即不同lane在同一寄存器编号下可以有相同的变量名，但实际上保存不同的值）。

此外，同一个wavefront内有共享的数据，它们在整个wavefront范围内是共用的，比如数组的基址。这类数据叫做标量状态，是通过 标量通用寄存器（Scalar General-Purpose Register，SGPR） 实现的。同一个wavefront 下的不同线程可以基于数组基址，再结合通道偏移量来找到自己要存取的数组数据。

0.3.2.2 局部数据共享存储与数据缓存

对于同一个工作组的不同线程（不管是不是属于同一个wavefront），如果要共享同步数据，AMD也在硬件上专门开辟了一块区域。这块区域叫做 局部数据共享存储（Local Data Share，LDS） ，在 HIP 中通过 __shared__ 进行声明。

前面提过，程序当然是能够访问全局内存的，但刚刚访问过的/经常访问的某全局内存数据会被放到缓存（Cache）中，方便后续快速访问（系统认为用户接下来需要继续使用），而不必再重新从全局内存中读取。缓存（Cache） 是由硬件进行管理的，如果某请求的数据能够在缓存中找到，则称为一次命中（Hit）。

下图是 VGPR、SGPR、LDS与缓存的示意图与相关使用过程：

0.3.3 向“合并访存”靠近

在GPU显存中，数据是按数据块进行组织与访问。一个数据块的大小一般来说是相对稳定的，就相当于格式化一块磁盘会将选定其数据块大小一样，所以说每个数据块中能够保存的数据是有限的。

这样的机制下，想要访问到某个数据，其实是先访问数据所在的整个数据块，然后再从数据块中取出我们想要的某个（些）数据。试想如果我们真正想要的数据分散在不同的数据块就会造成如下现象：先找到并传输所有相关的多个数据块，然后这些数据块只分别取出其中某个或某几个数据，这样造成了带宽的很大浪费，也增加访问的时间与效率。

显存的带宽是极其宝贵的，以 64 个工人（Lane）搬取货物（读取数据）为例说明合并访存的必要性：

合并访存成功时：64 个工人需要的东西都在同一个箱子里，卡车送来 1 个箱子，所有人同时拿到东西（合并访存，效率极高）。==【64 个线程的请求落入同一个数据块，GPU 只需要向显存（VRAM）发起 1 次 内存事务（Memory Transaction）就能把数据一次性搬运过来。】==
完全散乱时：64 个工人每人要的东西都在不同的仓库里，卡车必须跑 64 趟去搬 64 个箱子，而每个箱子装满数据，工人却只拿走了其中 4 个字节，造成严重的带宽浪费。==【GPU 的内存控制器不得不向显存发起 64 次 独立的内存事务。】==

程序员该如何解决这个问题？（数据结构与访问模式优化）既然待读取的数据位置决定了合并访存的成败，那么在 GPU 编程中，程序员的责任就是主动把数据“整理”成便于相邻线程连续读取的格式。最典型的优化手段就是 SOA 与 AOS 的转换：

例子：结构体数组（AOS）vs 数组结构体（SOA）假设要处理 1000 个粒子，每个粒子有 $(x, y, z)$ 坐标：

AOS (Array of Structures，结构体数组，对 GPU 很不友好)：
1
2
struct Particle { float x, y, z; }; Particle particles[1000];
如果每个线程负责一个粒子并读取其 x 坐标：
Lane 0 读 particles[0].x (地址 0)
Lane 1 读 particles[1].x (地址 12，中间隔了 y 和 z)
Lane 2 读 particles[2].x (地址 24) 结果：相邻线程访问的内存地址跨度大（跨步访问），数据离散，无法很好地合并访存。
SOA (Structure of Arrays，数组结构体，GPU 推荐做法)：
1
2
3
4
5
struct Particles { float x[1000]; float y[1000]; float z[1000]; };
如果每个线程读取 x 坐标：
Lane 0 读 x[0] (地址 0)
Lane 1 读 x[1] (地址 4)
Lane 2 读 x[2] (地址 8) 结果：数据在物理内存中是完全连续排列的！Lane 0~64 访问的恰好是连续的内存块，硬件能将其瞬间合并为一次连续的内存请求。

为此，在GPU编程中，不仅仅是要“写 Kernel 函数”，更核心的功夫在于设计数据在显存中的排布方式（尽量连续与对齐），以及合理分配线程的工作任务（让相邻线程恰好去读连续的数据），从而把分散的访存请求压缩成尽量少的数据块传输。

当同一个wavefront中不同线程访问数据时，我们应该尽量保证这些数据存储在相同/相邻的数据块。把同一 wavefront 的地址请求组织得便于硬件合并处理，这就是 合并访存（Memory Coalescing） 要解决的问题。当同一个wavefront中不同线程待访问的数据紧凑地分布在相同或相邻的数据块中时，硬件就能将其打包为极少数的内存事务（Memory Transactions），这就是合并访存（Memory Coalescing）。这将极大提升带宽传输、访存效率。

如下代码是 global_memory_access.hip 中关键片段：

template <unsigned Stride>
__global__ void gather_copy(
    const float* input, float* output, std::size_t n
) {
    std::size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
	    // (tid * Stride) & (n - 1) 的效果赞同于 (tid * Stride) % n
        std::size_t source = (tid * Stride) & (n - 1);
        output[tid] = input[source];
    }
}

代码中增加了一个变量Stride，通过给其赋予不同大小的值以控制与模拟待访问的数据是否在连续或相邻的位置。因为n 是 2 的幂，所以 (tid * Stride) & (n - 1) 的效果赞同于 (tid * Stride) % n，效果是把下标source控制在[0, n-1] 这个范围内。

读取步长Stride	时间中位数（ms）	逻辑有效带宽（GB/s）
1	0.235	571
17	0.772	174
257	1.89	71.1

其中“逻辑有效带宽”描述的是程序每秒完成了多少有用的读写工作。实验中每个元素读一次、写一次，所以逻辑有效带宽按 2 × N × sizeof(float) / 时间 计算。

这个表格展示了某次实验的结果数据：通过给Stride赋予不同的读取步长，模拟与计算了合并访存与非合并访存情况下相关结果参数的结果。

以下是我在设备"AMD Radeon 8060S Graphics"上测试得到的数据，好像跟上述表格中结果有点不一致：

(part0-intro) jovyan@jupyter-sanyinjiang:~/hello-gpu/code/part0-intro/chapter3$ ./global_memory_access --implementation all --size 16777216 --warmup 10 --repeat 50
ENV runtime=hip gpu="AMD Radeon 8060S Graphics" arch=gfx1151 hip_multiprocessor_count=20 wave_size=32
RESULT experiment=global-memory implementation=stride-1 runtime=hip shape=N=16777216 dtype=fp32 warmup=10 repeat=50 seed=20260726 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.572992 median_ms=0.574154 mean_ms=0.575007 logical_bandwidth_gbs=233.766069 tflops=0 max_abs_error=0
RESULT experiment=global-memory implementation=stride-17 runtime=hip shape=N=16777216 dtype=fp32 warmup=10 repeat=50 seed=20260726 timed=1 correct=OK precheck=OK postcheck=OK min_ms=5.152326 median_ms=5.185188 mean_ms=5.184238 logical_bandwidth_gbs=25.884832 tflops=0 max_abs_error=0
RESULT experiment=global-memory implementation=stride-257 runtime=hip shape=N=16777216 dtype=fp32 warmup=10 repeat=50 seed=20260726 timed=1 correct=OK precheck=OK postcheck=OK min_ms=2.177036 median_ms=2.181625 mean_ms=2.181773 logical_bandwidth_gbs=61.521909 tflops=0 max_abs_error=0

0.3.4 LDS的作用与工作流程

同一个工作组的多个不同的线程需要交换与复用某个（些）数据，就需要用到局部数据共享存储（LDS）。其中的多个线程按照约定流程将自己负责的那部分内容写进LDS，并在写入后更新 屏障同步（Barrier） 标识，以告诉其他线程“我已经将我负责的部分写入更新了，你们可以往下执行流程了”。当所有相关线程都写入自己负责的部分并更新 屏障同步（Barrier） 标识，待屏障解除后，各个线程再按照算法逻辑，安全地读取 LDS 中由自己或其他线程写入的数据进行后续计算。

通过这种“协同加载 + LDS 共享 + 屏障同步”的模式，原本需要多次访问慢速显存的重复操作，被转化为对高速 LDS 的直接访问，从而大幅提升了计算效率。

简单的数据相加、未重复读写数据的操作不需要使用LDS。

示例文件中 lds_bank_conflict.hip 中，工作组内的每个线程都会执行如下更新LDS数据、同步操作。

// 256 个线程合作初始化 8,448 个共享元素：每个线程以 block 大小为步长写多个位置。
// 所有 wavefront 的区域具有相同局部内容，因为 shared_value 对 index 按 kRegion 取模。
for (unsigned int offset = threadIdx.x; offset < kSharedElements;
     offset += kBlockSize) {
    shared[offset] = shared_value(offset);
}
// 确保所有共享内存写入对 block 内所有线程可见，之后才能开始读取。
__syncthreads();

写这类代码时，不能让一部分组内线程绕过一个其他线程需要到达的屏障。把同步放在后面的越界退出之前，也是先让组内线程完成共享数组初始化。屏障不负责修复越界、覆盖写入等其他错误，数组位置仍要分配正确。

0.3.5 wavefront的调度与执行

0.3.5.1 wavefront 三种状态：驻留、就绪和执行

由此引发对wavefront的几种状态的探讨：

状态	可以怎样理解
驻留（Resident）	工作已经获得相应片上资源，但因为数据等未准备就绪无法被立刻调度以执行。但其状态保留在硬件中
就绪（Ready / Eligible）	满足继续执行所需的条件，可以被选择
执行（Executing）	当前获得了执行机会，正在发射或执行相应指令

在整个系统中，可能会同时存在多个未完成的wavefront，硬件中可能保留其他wavefront的状态，如果当前wavefront因为数据未准备就绪从“执行”状态变更为“驻留”，那么其他处于”就绪“状态的wavefront就可能被调度到以执行。一般来说，系统总是尽可能让所有资源处于不停歇地工作中。

下图展示了系统中两个wavefront在未结束前具体状态的变更：延迟隐藏

上图中，wavefront A等数据需要一个过程，在这段时间内，wavefront A其实是相当于被卡着、等待数据的返回。期间，计算资源并没有闲着，它的使用权被交给wavefront B，保证系统中其他wavefront继续正常执行，从外面看来，整个系统还是处于工作状态，好像wavefront A等待数据期间的延迟被隐藏了一样，就个现象叫做延迟隐藏（Latency Hiding）。

有时存在极端的情况，即如果wavefront等待数据期间，系统中不存在处于“就绪”状态wavefront，那么系统中的计算资源就要空闲下来了。

0.3.5.2 wavefront 占用率

每份工作从开始执行到未结束期间，系统要保留其它的临时值、执行状态等，这些东西都要要占用GPU内部有限的片上资源（如寄存器、LDS等）。所以说每份使用与临时保留的资源要恰到好处，不能太多，也不能太少。太多会导致GPU上能够同时容纳的工作份数减少，太少的话可能造成频繁的读取、降低工作间共享数据的可能性。

以多个wavefront来说，GPU上的片上资源上限是固定的，如果单个wavefront的临时资源越多，则GPU上能够容纳的wavefront个数就越少。实际工作中还是考虑到分配粒度与其他上限，实际上能够容纳的wavefront个数一般比“GPU总片上资源数/wavefront平均占用临时资源数”小。

在某个GPU任务中，可以计算出一个大致可用的wavefront 槽位（个数上限）。处于“驻留”状态的 wavefront 个数 除以 可用的wavefront 槽位，再乘以100%，就是wavefront 占用率。

\[\text{wavefront 占用率} = \frac{\text{驻留 wavefront 个数}}{\text{可用的wavefront 槽位}} ✖️ \text{100\%}\]

当系统中 wavefront 占用率高时，平均来说处于就绪状态的wavefront量也相对高，此时可供选择的wavefront相对多，理论上计算资源可能会处于忙碌状态、提升其利用率，但实际工作中要警惕它们等待同一或同一类受限数据，此时可能造成计算资源处于饥饿状态、程序执行时间增加。

相对地，增加单个wavefront保留的临时资源，会使用系统中容纳wavefront数减少，从而减少wavefront占用率，就绪状态的wavefront可能减少了，但却会增加 wavefront 复用数据、内合并访存的可能性，从而提升程序执行效率。

所以是提升还是减少wavefron占用率，不能一概而论，要根据实际的内核设计与程序执行效果来分析与优化。

自我检验

线程要计算 c[2]，是否必须先把 a[2]、b[2] 放到 LDS？
两种实现里，每个 lane 都读取一个数，是否就能认为它们的访存效率一样？
一个线程把自己的共享位置写完，是否可以立即假设其他线程负责的位置也能读取？
GPU 上已经驻留很多 wavefront，是否代表总能找到就绪的工作？

回答：

不必。简单的数据相加、未重复读写数据的操作不需要使用LDS。
不一样。同一个wavefront中不同lane，如果是通过缓存、合并访存方式读取的数据效率更高。
不能。某个线程自己一般只对自己负责，无法绝对保证其他线程的进度。
不能保证。比如说它们等待同一或同一类受限数据，那么在获取到数据之前所以wavefront都仍是在驻留状态，不会进入就绪状态。

0.4 第一个程序 + 性能分析

本章节学习的主要内容包括：准备或恢复第1章那个软硬件环境；对code/part0-intro/chapter4/vector_add.hip这个HIP程序关键代码进行解释并编译执行它；学习通过PyTorch构架在CPU与AMD GPU上执行并行计算，规定了统计口径的情况下，重点统计加法GPU执行流中核函数本身在设备侧的整个执行时间、每次统计取平均值/中位数/最小值来对比有效数据带宽及与官方最高带宽进行对比；因为受到GPU 侧排队、调度或与该操作相关的其他设备端开销等因素影响，当前方法得到的 elapsed_time 并不能严格等同于“加法 kernel 的执行时间”，进行多次测试综合统计可能能得出较可信的结果；给出了不同统计口径下，锚点的添加位置不同。

0.4.1 准备环境

略。见自己前面的记录：环境准备

0.4.2 从向量加法到 HIP 程序

这部分主要是对code/part0-intro/chapter4/vector_add.hip这个HIP程序关键代码进行解释，自己有C++/C语言基础，掌握基本的常用语法、容器使用，然后自己曾经阅读一些CUDA C的语法基础入门，最后还有开源教程中提供了关键解释，所以此示例阅读起来没有问题。

相当于是自己回顾了一下以前学习过的C++/C语言基础、CUDA C的语法基础。

最后编译、执行是简单的命令操作。

(hello-gpu-notebook-part0-intro) root@j-7qhdfxoi6o3pyecodbv33kwd7z:/workspace/notebooks/part0-intro/chapter4# hipcc vector_add.hip -O2 -o vector_add
(hello-gpu-notebook-part0-intro) root@j-7qhdfxoi6o3pyecodbv33kwd7z:/workspace/notebooks/part0-intro/chapter4# ./vector_add 
device_name: AMD Radeon Pro W7900D
vector_size: 1048576
blocks: 4096
threads_per_block: 256
max_error: 0
status: PASS

### 0.4.3 测量 PyTorch 向量加法

其中使用PyTorch直接在AMD显卡上执行并行计算，只统计GPU上执行加法本身的时间开销，更具体地说是：

GPU 侧用同一执行流中的 start Event → torch.add → end Event 测量。
不包含输入创建、输出分配、CPU 与 GPU 间的数据拷贝、结果验证。
不包含 Python 调用、PyTorch 框架分派，以及 CPU 向 GPU 提交 kernel 的主机侧启动开销。
Event 得到的是设备时间线中这段操作的经过时间；它适合比较该 torch.add 实现，但不能严格等同于“纯 GPU 算术指令执行时间”。

PyTorch 中核函数启动仍然由 CPU 发起，只是被框架封装在这句里了：

1	`torch.add(a, b, out=c)`

可以把它理解为这条 Python 语句在 CPU 侧大致完成了：

Python 调用 torch.add
  → PyTorch 根据 a/b/c 的设备、类型、形状选择 ROCm 实现
  → 准备 kernel 参数
  → 调用 HIP 运行时，将 kernel 启动命令放入 GPU 的执行流
  → Python 返回

最后一步在 HIP 手写程序中会显式写成：

1	`vector_add<<<blocks, threads>>>(d_a, d_b, d_c, n);`

而 PyTorch 把类似的代码藏在其 C++/ROCm 后端内部，所以在 Python 脚本中看不到 <<<...>>> 或 hipLaunchKernelGGL(...)。本例的一轮 GPU 时间线可以简化为：

CPU：
start.record() → torch.add(...) → end.record() → end.synchronize()
                    ↑
             CPU 在这里完成框架分派与 kernel launch

GPU 执行流：
start Event → 加法 kernel → end Event

因此：

torch.add(...) 这行的 CPU 执行时间包含“发起 kernel”的开销。
start.elapsed_time(end) 测得的是 GPU 执行流中两个 Event 之间的设备侧时间。
CPU 发起 kernel 所花的 Python、PyTorch 分派和 HIP 提交时间，不在 GPU Event 测得的时间内。
end.synchronize() 会让 CPU 等待 GPU 完成，但它位于 Event 区间之后；其等待时间也不等于 GPU Event 的结果。

对于当前未使用 torch.compile 的简单 torch.add，通常可近似理解为：PyTorch 在内部启动一个逐元素加法 kernel。但这是框架实现细节，不应把 Event 结果严格称为“纯 GPU 算术指令执行时间”（Event 结果即start.elapsed_time(end)，其主要是加法 kernel 的执行时间，但可能还包含 GPU 侧排队、调度或与该操作相关的其他设备端开销；但不能把这些部分精确分离出来。）。

此外，作为参照，也在CPU上统计了 torch.add 调用前后的 CPU 时钟计时，包含 Python 调用、PyTorch 分派和底层 CPU 算子执行时间。

0.4.4 解释测量结果

PyTorch 向量加法示例中，每次加法需要处理的有效数据量为 3个float32数，即3 * 4 = 12 Bytes。假设示例需要处理的数组长度是N，那么整个示例算法需要处理的有效数据量为：12 * N 。

再设定 CPU/GPU 上执行计算的耗时是 t （单位：ms），那么得到有效带宽： \[ BW_{\mathrm{effective}}=\frac{12N}{t}. \]

有效带宽回答的是“每秒完成了多少算法所要求的数据处理”，并没有测出显存实际传输了多少字节。缓存、访存粒度和写入路径都可能让物理流量与不同。

上述公式再变更下形式与单位，得到如下计算公式（结果单位是：GB/s）： \[ BW_{\mathrm{effective}} =\frac{B_{\mathrm{effective}}\times10^{-6}}{耗时\_ms} \]

0.4.5 从算子时间到程序时间

此程序中涉及到的统计时间，要先给限定统计口径，不同的统计口径得到的结果不同，这是一个细致而相对复杂的过程。

想回答的问题	应观察的范围
数据已在设备上，一次 PyTorch 加法经过多久？	本章的 GPU event 区间
主机发出操作并等到结果完成，要多久？	包含提交与等待的主机计时
CPU 内存中的两个数组交给 GPU，直到取回结果，要多久？	加上主机与设备间拷贝；说明是否包含分配

0.4.6 记录与练习

0.4.6.1 回答问题

保持每块 256 个线程，若输入长度改成 1000，需要几个 block？最后一块中哪些线程会被 idx < n 排除？这只是覆盖范围的推导，不代表已经运行过修改后的 HIP 程序。
若开始和结束 event 仍按原顺序提交，却在读取 elapsed_time 前不等待，缺少了什么保证？
使用自己测得的 GPU 中位数，手算。结果是否与 gpu_bandwidth_gb_s_by_median 相符？先检查毫秒与秒的换算。
比较三个规模的 mean、median 和 min。它们是否足够接近？再测一组时是否仍得到相同趋势？记录观察，不预设输入越大就一定越有效率。

回答： 1. 需要的线程块数为 (1000 + 256 -1)/256，结果是4，即需要4个线程块。前三个线程块共有256 * 3= 768个线程，最后一个线程块还有（1000-768）=232个线程实际参与计算，所以最后一个线程块即编号为3（线程块编号从0开始）的线程块中，实际参与计算的块内线程编号是0~231，编号为 232~255 这些块内线程会被idx < n 排除。 2. 无法保证GPU上本轮工作流中所有任务已经完成。导致计算得到的耗时可能比实际的小，从而导致后续相关统计结果不正确。 3. 我先记录下自己实际执行时没得的数据： (1)python benchmark_vector_add.py --size 1048576 --warmup 5 --repeat 30 --cpu-threads 1

(hello-gpu-notebook-part0-intro) root@j-7qhdfxoi6o3pyecodbv33kwd7z:/workspace/notebooks/part0-intro/chapter4# python benchmark_vector_add.py --size 1048576 --warmup 5 --repeat 30 --cpu-threads 1
torch: 2.13.0+rocm10.0.0
hip: 7.15.26333
cuda_available: True
device_name: AMD Radeon Pro W7900D
cpu_name: AMD EPYC 9334 32-Core Processor
cpu_threads: 1
dtype: float32
vector_size: 1048576
warmup: 5
repeat: 30
cpu_timing: wall clock around torch.add(a, b, out=c)
gpu_timing: device event interval; not end-to-end latency
cpu_validation: PASS (all elements finite and equal to 3)
gpu_validation: PASS (all elements finite and equal to 3)
cpu_mean_ms: 0.132728
cpu_median_ms: 0.131815
cpu_min_ms: 0.130460
cpu_bandwidth_gb_s_by_median: 95.459066
cpu_bandwidth_gb_s_by_min: 96.450592
gpu_mean_ms: 0.030736
gpu_median_ms: 0.029359
gpu_min_ms: 0.024360
gpu_bandwidth_gb_s_by_median: 428.580610
gpu_bandwidth_gb_s_by_min: 516.539918
status: PASS

根据计算公式： \[ BW_{\mathrm{effective}} =\frac{B_{\mathrm{effective}}}{gpu\_median\_ms\times10^{-3}} \] 可得 $$ BW_{} = ^{11}

=428 . $$

(2)python benchmark_vector_add.py --size 16777216 --warmup 5 --repeat 30 --cpu-threads 1

(hello-gpu-notebook-part0-intro) root@j-7qhdfxoi6o3pyecodbv33kwd7z:/workspace/notebooks/part0-intro/chapter4# python benchmark_vector_add.py --size 16777216 --warmup 5 --repeat 30 --cpu-threads 1
torch: 2.13.0+rocm10.0.0
hip: 7.15.26333
cuda_available: True
device_name: AMD Radeon Pro W7900D
cpu_name: AMD EPYC 9334 32-Core Processor
cpu_threads: 1
dtype: float32
vector_size: 16777216
warmup: 5
repeat: 30
cpu_timing: wall clock around torch.add(a, b, out=c)
gpu_timing: device event interval; not end-to-end latency
cpu_validation: PASS (all elements finite and equal to 3)
gpu_validation: PASS (all elements finite and equal to 3)
cpu_mean_ms: 6.386754
cpu_median_ms: 6.300809
cpu_min_ms: 6.250024
cpu_bandwidth_gb_s_by_median: 31.952499
cpu_bandwidth_gb_s_by_min: 32.212132
gpu_mean_ms: 0.299717
gpu_median_ms: 0.289359
gpu_min_ms: 0.288039
gpu_bandwidth_gb_s_by_median: 695.766323
gpu_bandwidth_gb_s_by_min: 698.956019
status: PASS

根据公式可得 $$ BW_{} = ^{11}

=695 . $$

(3)python benchmark_vector_add.py --size 67108864 --warmup 5 --repeat 30 --cpu-threads 1

(hello-gpu-notebook-part0-intro) root@j-7qhdfxoi6o3pyecodbv33kwd7z:/workspace/notebooks/part0-intro/chapter4# python benchmark_vector_add.py --size 67108864 --warmup 5 --repeat 30 --cpu-threads 1
torch: 2.13.0+rocm10.0.0
hip: 7.15.26333
cuda_available: True
device_name: AMD Radeon Pro W7900D
cpu_name: AMD EPYC 9334 32-Core Processor
cpu_threads: 1
dtype: float32
vector_size: 67108864
warmup: 5
repeat: 30
cpu_timing: wall clock around torch.add(a, b, out=c)
gpu_timing: device event interval; not end-to-end latency
cpu_validation: PASS (all elements finite and equal to 3)
gpu_validation: PASS (all elements finite and equal to 3)
cpu_mean_ms: 21.338049
cpu_median_ms: 21.250562
cpu_min_ms: 21.035420
cpu_bandwidth_gb_s_by_median: 37.895768
cpu_bandwidth_gb_s_by_min: 38.283351
gpu_mean_ms: 1.121033
gpu_median_ms: 1.121419
gpu_min_ms: 1.118158
gpu_bandwidth_gb_s_by_median: 718.114016
gpu_bandwidth_gb_s_by_min: 720.208039
status: PASS

根据公式可得 $$ BW_{} = ^{11}

=718 . $$ 三个手算结果，分别与对应的程序输出结果是一致的。

将它们绘制成一张表格如下：三组均于 2026-10-23 在上述 Radeon Pro W7900D / 原生 Ubuntu 24.04.4 环境实测，以下只比较 PyTorch GPU 路径，均为 float32、warmup 5、repeat 30：

元素数	每个数组的大小	GPU 平均数/中位数/最小数（单位：ms）	按平均数/中位数/最小数计算的有效带宽（单位：GB/s）
1,048,576	4 MiB	0.030736 / 0.029359 / 0.024360	409.386778 / 428.587895 / 516.539901
16,777,216	64 MiB	0.299717 / 0.289359 / 0.288039	671.722298 / 695.767514 / 698.956016
67,108,864	256 MiB	1.121033 / 1.121419 / 1.118158	718.360983 / 718.113718 / 720.208028

不完全接近，且规模越大越接近、越稳定。大致规律是：

小规模（4 MiB）中，加法本身极短，GPU 提交、调度、时钟状态和其他进程活动等固定或偶发因素占比更高，所以 mean、median、min 差异较大。
中等规模（64 MiB）开始，实际数据处理时间占主要部分；min 与 median 已很接近，但少数慢轮次仍会使 mean 偏大。
大规模（256 MiB）时，三种统计量几乎一致，说明该配置下测量较稳定。

带宽也呈现对应趋势：按中位数计算，从约 428.6 GB/s 增至 695.8 GB/s、718.1 GB/s。这说明随着规模增大，固定开销被更多数据处理工作摊薄，性能逐步接近该实验环境中的稳定带宽区间。根据 AMD Radeon专业显示参数可知，AMD 为 Radeon Pro W7900D 给出的规格是最高 864 GB/s 显存带宽。$720/864\approx83.3\%$ 可以作为模型下的对照，但不能称为显存控制器的实测利用率，也不能据此判断“只剩 16.7% 优化空间”。

0.4.6.2 报错处理

0.4.6.2.1 提示`orch.AcceleratorError: CUDA error: device kernel image is invalid`

尝试在容器中执行如下命令时，

(hello-gpu-notebook-part0-intro) root@j-7qhdfxoi6o3pyecodbv33kwd7z:/workspace/notebooks/part0-intro/chapter4# python benchmark_vector_add.py --size 1048576 --warmup 5 --repeat 30 --cpu-threads 1
torch: 2.13.0+rocm10.0.0
hip: 7.15.26333
cuda_available: True
device_name: AMD Radeon Pro W7900D
cpu_name: AMD EPYC 9334 32-Core Processor
cpu_threads: 1
dtype: float32
vector_size: 1048576
warmup: 5
repeat: 30
cpu_timing: wall clock around torch.add(a, b, out=c)
gpu_timing: device event interval; not end-to-end latency
Traceback (most recent call last):
  File "/workspace/notebooks/part0-intro/chapter4/benchmark_vector_add.py", line 173, in <module>
    main()
  File "/workspace/notebooks/part0-intro/chapter4/benchmark_vector_add.py", line 161, in main
    gpu_times = benchmark_gpu(args.size, args.warmup, args.repeat)
                ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/workspace/notebooks/part0-intro/chapter4/benchmark_vector_add.py", line 79, in benchmark_gpu
    a = torch.ones(size, device="cuda", dtype=torch.float32)
        ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
torch.AcceleratorError: CUDA error: device kernel image is invalid
Search for `hipErrorInvalidImage' in https://rocm.docs.amd.com/projects/HIP/en/latest/index.html for more information.
For more detailed error information, run with CUDA_LOG_FILE=stderr
Device-side assertion tracking was not enabled by user.
(hello-gpu-notebook-part0-intro) root@j-7qhdfxoi6o3pyecodbv33kwd7z:/workspace/notebooks/part0-intro/chapter4#

原因高度确定：容器中的 PyTorch 很可能安装了默认的 gfx1201（RX 9070 XT）设备包，但租赁卡 Radeon Pro W7900D 属于 gfx1100。

cuda_available: True 只说明 ROCm 能看见 GPU；直到这句首次真正提交 GPU 算子时：

1	`a = torch.ones(size, device="cuda", dtype=torch.float32)`

PyTorch 才加载并执行 GPU 内核。若内核镜像只含 gfx1201、不含 gfx1100，就会出现：

1 2	`hipErrorInvalidImage CUDA error: device kernel image is invalid`

这不是 benchmark_vector_add.py 的逻辑问题，也与 c = torch.empty_like(a) 无关。torch: 2.13.0+rocm10.0.0 和 hip: 7.15.26333 的组合本身是正常的。

先在容器中确认实际架构和安装的设备包：

python - <<'PY'
import importlib.metadata
import torch

prop = torch.cuda.get_device_properties(0)
print("GPU:", prop.name)
print("architecture:", prop.gcnArchName)

for dist in importlib.metadata.distributions():
    name = dist.metadata["Name"]
    if name and ("gfx" in name.lower() or name.lower().startswith("amd-torch")):
        print(name, dist.version)
PY

执行时能看到：

1 2	`GPU: AMD Radeon Pro W7900D architecture: gfx1100`

若输出中出现 amd-torch-device-gfx1201、rocm-sdk-device-gfx1201，却没有对应的 gfx1100 包，就证实了上述原因。

修复方式是在 part0-intro/pyproject.toml 中将以下三处统一从 device-gfx1201 改为 device-gfx1100：

1
2
3

"torch[device-gfx1100]==2.13.0+rocm10.0.0",
"torchvision[device-gfx1100]==0.28.0+rocm10.0.0",
"rocm[devel,device-gfx1100]==10.0.0",

# 以下是修复后的部分内容：
dependencies = [
    #"torch==2.13.0+rocm10.0.0",
    "rocm[devel]==10.0.0",
    "hello-gpu-notebooks",
    "ipykernel>=7.4.0",
    "jupyterlab>=4.6.4",
    "nbclient>=0.11.0",
    "nbformat>=5.11.1",
    "pytest>=9.1.1",
    "triton>=3.8.0",
    "numpy>=2.5.3",
    "torch[device-gfx1100]==2.13.0+rocm10.0.0",
    "torchvision[device-gfx1100]==0.28.0+rocm10.0.0",
    "rocm[devel,device-gfx1100]==10.0.0",
]

然后在 part0-intro 目录执行：

1
2
3

uv sync
source ./activate-rocm.sh
python chapter1/check_torch_rocm.py

(hello-gpu-notebook-part0-intro) root@j-7qhdfxoi6o3pyecodbv33kwd7z:/workspace/notebooks/part0-intro# python chapter1/check_torch_rocm.py
python: 3.12.3
torch: 2.13.0+rocm10.0.0
cuda_available: True
device_count: 1
device_name: AMD Radeon Pro W7900D
result_shape: (1024, 1024)
result_dtype: torch.float32
result_device: cuda:0
result_checksum: 2210.884766

可以看到第 1 章的 PyTorch GPU smoke test 通过，再执行上述程序成功。

0.4.7 本章小结

一个完整 HIP 示例包含主机数据准备、设备存储、数据拷贝、核函数启动、等待和结果检查；核函数只是其中一部分。
一线程处理一元素时，全局下标决定访问位置，边界判断保护最后一个 block 的无效线程。
手写 HIP 程序负责验证结果，PyTorch benchmark 负责演示计时。测量前需要明确实现、输入和计时范围。
GPU event 标记区间，同步保证区间完成；预热后重复测量，用中位数配合平均值和最小值阅读结果。
给出向量加法的有效带宽。它是算法数据量与实测时间的比值，不能代替物理显存流量或端到端程序耗时。

算子优化与Agent开发

#GPU #算子 #Agent

HelloGPU-P00-入门与硬件速通

https://jiangsanyin.github.io/2026/09/20/HelloGPU-P00-入门与硬件速通/

作者

sanyinjiang

发布于

2026年9月20日

许可协议

HelloGPU-P01-Profiling实战上一篇

Hermes部署与使用下一篇

HelloGPU-P00-入门与硬件速通

第0篇 入门与硬件速通

0.0 项目介绍与学习提示

0.1 环境准备

0.1.1 使用AUP Learning Cloud环境

根据环境修改配置

同步本篇 uv 环境：

验证 GPU 可见性

验证 PyTorch ROCm

验证最小 HIP 程序

0.1.2 使用AMD Radeon Cloud环境（优先推荐）

根据环境修改配置

同步本篇 uv 环境：

验证 GPU 可见性

验证 PyTorch ROCm

验证最小 HIP 程序

0.2 GPU 体系结构（上）：编程模型与 wavefront 执行

0.2.1 GPU计算的基本流程

0.2.2 线程的组织与执行

线程块与网格

线程块内的 wavefront

0.2.3 线程编号与数据索引

0.2.3.1 线程块与块内线程的编号

0.2.3.2 数组下标的计算

0.2.3.3 数组边界检查

0.2.3.4 线程编号与lane编号

0.2.3.5 线程分组与硬件执行单元

0.2.4 本章总结

0.2.5 自我检验

0.2.6 完整做vector_add的HIP示例

0.3 GPU 体系结构（下）：片上资源与数据通路

0.3.1 数据位置与读写过程

0.3.2 Wavefront 的访存组织

0.3.3 Block 内的数据共享与同步

0.3.4 访存等待与指令调度

本章小结

自我检验

0.3 GPU 体系结构（下）：片上资源与数据通路（旧）

0.3.1 一次加法的大致过程

0.3.2 数据的存取、管理与缓存

0.3.2.1 向量通用寄存器与标量通用寄存器

0.3.2.2 局部数据共享存储与数据缓存

0.3.3 向“合并访存”靠近

0.3.4 LDS的作用与工作流程

0.3.5 wavefront的调度与执行

0.3.5.1 wavefront 三种状态：驻留、就绪和执行

0.3.5.2 wavefront 占用率

自我检验

0.4 第一个程序 + 性能分析

0.4.1 准备环境

0.4.2 从向量加法到 HIP 程序

0.4.4 解释测量结果

0.4.5 从算子时间到程序时间

0.4.6 记录与练习

0.4.6.1 回答问题

0.4.6.2 报错处理

0.4.6.2.1 提示orch.AcceleratorError: CUDA error: device kernel image is invalid

0.4.7 本章小结

第0篇入门与硬件速通

0.4.6.2.1 提示`orch.AcceleratorError: CUDA error: device kernel image is invalid`