AMDGPU驱动（准确地说是ROCm软件栈，特别是HIP runtime和HSA runtime）解析hipcc编译的可执行文件的过程，是一个从ELF段读取到用户态Runtime解析，最后通过内核驱动加载的多层协作过程。

简单来说，hipcc将GPU代码（HSACO）打成由特定格式包裹的”Fat Binary”，并嵌入到Host程序的ELF文件中。程序运行时，通过自动生成的构造函数将这段数据的指针注册给HIP Runtime，Runtime再根据当前GPU型号提取对应的代码并加载。

以下是详细的解析流程：

1. 编译阶段：Fat Binary 的生成与嵌入

在理解解析之前，必须先知道文件结构是怎么来的。hipcc（实际调用clang）的编译流程如下：

• 设备侧编译: 源代码被编译为AMD GPU的机器码，即HSA Code Object (HSACO)。如果你指定了多个架构（如 –offload-arch=gfx90a –offload-arch=gfx942），会生成多个HSACO文件。

• 打包 (Bundling): 使用 clang-offload-bundler 工具，将这些不同架构的HSACO文件打包成一个 Clang Offload Bundle。这是一个包含Magic Number（通常是CCOB）和目录索引的二进制格式。

• 嵌入 (Embedding): 这个Bundle被作为数据段嵌入到最终的Host ELF可执行文件中。
  Section Name: .hip_fatbin
  Symbol: __hip_fatbin (这是一个全局符号，指向该段的起始地址

2. 启动阶段：自动注册 (Self-Registration)

当你运行可执行文件时，操作系统加载器（Loader）将ELF文件映射到内存。此时GPU驱动并未介入，而是依靠编译器插入的Host侧构造函数来触发解析：

• 构造函数: 编译器在每个编译单元中插入了一个全局构造函数（类似于 attribute((constructor))）。

• 注册调用: 这个构造函数会执行 __hipRegisterFatBinary(&__hip_fatbin)。

  • __hip_fatbin 就是指向内存中 .hip_fatbin 段的指针。
  • 该函数是 libamdhip64.so (HIP Runtime) 提供的API

3. HIP Runtime 解析逻辑

当 __hipRegisterFatBinary 被调用时，HIP Runtime (用户态驱动) 开始工作。它并不会立即把所有代码加载到GPU，而是建立索引:

• 识别 Bundle 格式: Runtime 检查传入指针的头部，识别 Clang Offload Bundle 格式（检查魔数 CCOB 或旧版的 __CLANG_OFFLOAD_BUNDLE__）。
• 解析目录 (Index Parsing): Runtime 解析Bundle的头部信息，读取其中包含的所有 Bundle Entry ID。
  • Entry ID 格式通常为：hip-amdgcn-amd-amdhsa–gfx90a (指明了语言-架构-厂商-OS-具体GPU型号)。
  • 匹配设备: Runtime 此时会查询系统中的GPU设备信息（通过 rocminfo 或 HSA 接口获取当前显卡的 GFX ID，例如 gfx90a）。
  • 提取代码 (Extraction):
  • Runtime 遍历 Bundle 目录，寻找与当前系统GPU架构最匹配的 Entry。
  • 一旦找到，计算出该 HSACO 在内存中的偏移量和大小，将其“解包”（逻辑上提取，实际是指向该内存片段.

4. 加载到 GPU (HSA Runtime & KMD)

拿到提取出来的 HSACO (ELF格式的GPU二进制) 后，HIP Runtime 将其交给更底层的 HSA Runtime (ROCr) 进行加载：

• 创建可执行对象: 调用 hsa_executable_create 创建一个空的GPU可执行容器。
• 加载代码对象: 调用 hsa_executable_load_code_object。
  • HSA Runtime 解析 HSACO 的 ELF 格式，处理其中的 .text (指令)、.data (全局变量) 等段。
  • 它会解析 Kernel Descriptor，这是AMD GPU特有的元数据，包含寄存器使用量（SGPR/VGPR）、LDS大小等信息。

5. 内核态交互 (KMD)

HSA Runtime 通过 ioctl 系统调用与内核驱动 amdgpu (KFD部分) 通信。
驱动负责分配显存，将代码段和数据段DMA到GPU显存中，并设置页表。

总结流程图

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29

[ELF Executable File]
      |
      +-- Section: .hip_fatbin (包含了所有架构的GPU代码)
               |
[程序启动 (Host CPU)]
      |
      +-- Constructor Function (__hip_fatbin_ctor)
               |
               v
      +-- Call __hipRegisterFatBinary(&__hip_fatbin)
               |
[HIP Runtime (User Space)] <--- 此时开始解析
      |
      +-- 1. 识别 Clang Offload Bundle (Magic: CCOB)
      +-- 2. 遍历 Bundle Entry IDs
      +-- 3. 匹配当前 GPU (例如: 找到 gfx90a)
      +-- 4. 提取对应的 HSACO 二进制数据
               |
               v
[HSA Runtime (ROCr)]
      |
      +-- 解析 HSACO ELF (Kernel Descriptors, Symbols)
      +-- hsa_executable_load_code_object
               |
               v
[AMDGPU Kernel Driver (KFD)]
      |
      +-- 分配显存，将代码上传至 VRAM
      +-- 准备好 Kernel Handle 供后续 hipLaunchKernel 使用

引言

这是一篇用来测试博客是否渲染成功的硬核技术文章。

核心代码测试

1
2
3
4

#include <hip/hip_runtime.h>
__global__ void hello() {
    printf("Hello from GPU!\n");
}

引言

AMD开源了runtime library, 如果你有一张AMD的卡， 完全可以自己尝试编译一下他们的仓库，debug一下

步骤

克隆仓库: https://github.com/ROCm/rocm-systems.git和https://github.com/ROCm/llvm-project.git, 我自己编译的版本是rocm-7.2, 注意分支的一致性

编译LLVM

下面的路径可以按需替换

1
2
3
4
5
6
7
8
9
10

cmake -S llvm -B build -DCMAKE_BUILD_TYPE=Release -G Ninja \
      -DLLVM_ENABLE_PROJECTS="clang;lld" \
      -DLLVM_EXTERNAL_PROJECTS="device-libs;comgr" \
      -DLLVM_EXTERNAL_DEVICE_LIBS_SOURCE_DIR=/home/jing.zhou/rocm-llvm/amd/device-libs \
      -DLLVM_EXTERNAL_COMGR_SOURCE_DIR=/home/jing.zhou/rocm-llvm/amd/comgr \
      -DCMAKE_INSTALL_PREFIX=/home/jing.zhou/rocm-systems/projects/install \
      -DLLVM_TARGETS_TO_BUILD="X86;AMDGPU" \
      -DCMAKE_LINKER=/home/jing.zhou/tools/bin/mold

编译clr

cd到rocm-systems目录，

1
2
3
4
5
6
7
8
9
10
11

export CLR_DIR="$(readlink -f projects/clr)"
export HIP_DIR="$(readlink -f projects/hip)"

export PATH=/home/jing.zhou/rocm-systems/projects/install/bin:$PATH

cd "$CLR_DIR"
mkdir -p build; cd build
cmake -DHIP_COMMON_DIR=$HIP_DIR -DCMAKE_PREFIX_PATH="/home/jing.zhou/rocm-systems/projects" -DROCM_INSTALL_PATH="/home/jing.zhou/rocm-systems/projects/install" -DCMAKE_INSTALL_PREFIX=/home/jing.zhou/rocm-systems/projects/install -DCLR_BUILD_HIP=ON -DCLR_BUILD_OCL=OFF -DHIP_PLATFORM=amd ..
make -j$(nproc)
make install

demo

1
2
3
4

export LD_LIBRARY_PATH=/home/jing.zhou/rocm-systems/projects/install/lib:$LD_LIBRARY_PATH
export HIP_CLANG_PATH="/home/jing.zhou/rocm-systems/projects/install/bin"
export HIP_PATH="/home/jing.zhou/rocm-systems/projects/install"
export HIP_PLATFORM=amd

写一份简单的hello-world代码

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35

#include <iostream>
// 引入 HIP 运行时头文件
#include <hip/hip_runtime.h>

// 1. 定义 GPU 核函数 (Kernel)
// __global__ 告诉编译器这个函数是在 GPU 上执行的，但由 CPU 发起调用
__global__ void gpu_hello_world() {
  // blockIdx, blockDim, threadIdx 是 HIP 的内置变量
  // 这里计算出当前线程的全局唯一 ID
  int gtid = blockIdx.x * blockDim.x + threadIdx.x;

  // printf 在 GPU 核函数中是受支持的（适合调试）
  printf("Hello World from GPU thread %d!\n", gtid);
}

int main() {
  // 2. 在 CPU (Host) 上打印
  std::cout << "Hello World from CPU!" << std::endl;

  // 3. 配置 GPU 的执行网格 (Grid) 和线程块 (Block)
  // 这里我们启动 1 个线程块，里面包含 4 个线程
  dim3 grid(1);
  dim3 block(4);

  // 4. 启动 GPU 核函数
  // 使用 <<<...>>> 语法来指定网格和线程块大小
  hipLaunchKernelGGL(gpu_hello_world, grid, block, 0, 0);

  // 5. 同步 CPU 和 GPU
  // 因为 GPU 的执行是异步的，CPU 需要等待 GPU 打印完毕后再退出程序
  hipDeviceSynchronize();

  return 0;
}

输出

1
2
3
4
5
6

Hello World from CPU!
Hello World from GPU thread 0!
Hello World from GPU thread 1!
Hello World from GPU thread 2!
Hello World from GPU thread 3!