讲 ROCm 各层功能和调用关系的文章不少，但真正抓住某一个概念、从上到下一条龙拆到底的却很少。因为工作恰好完整趟过整个栈，这里我把自己的梳理总结成一张对照表，分享给同样想搞懂 ROCm 分层的技术读友。

---

先抛个问题勾一下。

你在 PyTorch 或自己的 HIP 代码里写下这么两行：

hipMalloc(&ptr, n);              // 申请一块显存
hipLaunchKernelGGL(kernel, ...); // 启动一个核函数

它们看着简单，但每一行往下都要穿过 HIP → ROCr → libhsakmt → KFD 四层，才最终落到 GPU 硬件上。麻烦的是：同一个概念在每一层都换了个名字——你以为的"流（stream）“，到了下一层叫"队列（queue）”，再往下又变成内核里的一个 MQD(Memory Queue Descriptor)。名字对不上，读源码时就容易断片。

这里我试图把这套分层和"改名规律"彻底讲清楚，给你一张能贴在显示器边上的对照表。

1、先立分层模型：四层各管什么

ROCm 计算栈从上到下是四层，职责非常清晰：

一句话概括每层的活：

• HIP：给你一套顺手的、CUDA 风格的 API（hip*）。它的核心价值是可移植——同一份代码能在 AMD 和 N 卡上跑。但我觉得正是为了学CUDA，导致hip层概念向下对应时乱乱的，不知道兼容CUDA的国产GPU的软件栈是不是做的更好些，欢迎讨论。
• ROCr / HSA 运行时：把 HIP 的请求翻译成 HSA 标准里的对象——agent、queue、signal、AQL 包。这层是"标准层"，定义了异构计算该长什么样。
• libhsakmt（Thunk）：一层很薄的库，几乎不做逻辑，只把 HSA 的操作翻译成对内核 KFD 的 ioctl 调用。名字里的 kmt 就是 Kernel Mode Thunk。
• KFD：内核态计算驱动，真正去创建队列、分配 BO、注册中断、操作 GPUVM。

记住一个心法：越往上越"好用"，越往下越"贴硬件"；每下一层，都是把上一层的抽象翻译成更接近硬件的语言。

2、镇文之表：核心概念四层对照

这是全文最该收藏的一张表。把最常用的六类概念，在四层里的对应名字全列出来：

概念	HIP 层	ROCr / HSA 层	libhsakmt 层	KFD 内核
设备	hipDevice_t / hipCtx_t	hsa_agent_t（GPU agent）	HsaNodeProperties / NodeId	topology node、gpu_id
流 / 队列	hipStream_t	hsa_queue_t（AQL 队列）	hsaKmtCreateQueue	create_queue ioctl、MQD、doorbell
事件 / 信号	hipEvent_t	hsa_signal_t	hsaKmtCreateEvent / HsaEvent	kfd event、中断
kernel 启动	hipLaunchKernelGGL	AQL dispatch packet + 敲 doorbell	doorbell 映射	CP / MES 取包执行微码
显存分配	hipMalloc	hsa_amd_memory_pool_allocate	hsaKmtAllocMemory + MapMemoryToGPU	alloc BO ioctl、GPUVM map
统一内存 / SVM	hipMallocManaged（分配） hipMemAdvise（设访问偏好）	hsa_amd_svm_attributes_set / hsa_amd_svm_prefetch_async	hsaKmtSVMSetAttr	kfd svm ioctl、HMM 按需迁移

看这张表最该体会的，是改名背后的规律：不是随便换词，而是每层站的抽象高度不同——HIP 关心"我这个程序里的一条流"，HSA 关心"符合标准的一个队列对象"，Thunk 关心"该调哪个 ioctl"，KFD 关心"内核里那个 MQD 结构和 doorbell 页"。名字变了，是因为关注点变了。

下面挑三条最典型的"翻译链"走一遍，你就彻底通了。

3、翻译链一：设备是怎么从 hipDevice 变成 gpu_id 的

• 你用 hipSetDevice(0) 选卡，HIP 眼里它就是"第 0 号设备"。
• 到 ROCr，这块卡是一个 hsa_agent_t——注意 HSA 里 CPU 也是 agent，GPU 只是其中一种，这是标准化带来的统一抽象。
• 再往下，libhsakmt 通过 hsaKmtAcquireSystemProperties 拿到系统里所有拓扑节点（node），每个 GPU 对应一个 NodeId。
• 最终在 KFD 里，它是拓扑表中的一项，有个唯一的 gpu_id，你甚至能在 /sys/class/kfd/kfd/topology/ 下看到。

一句话：同一块卡，从"第几号"一路具体成"内核拓扑里的 gpu_id"。

4、翻译链二：stream 到底是不是 queue

这条最容易让人懵，因为 HIP 用户天天说"流"，但底层根本没有"流"这个东西——只有队列。

• HIP 的 hipStream_t 是"一串按顺序执行的任务"这个编程概念。
• 到 ROCr，它落地成一个 hsa_queue_t——一块用户态可写的环形缓冲区，你把 AQL 包（HSA 的标准命令格式）写进去。
• libhsakmt 用 hsaKmtCreateQueue 把这个队列注册到内核。
• KFD 在内核里为它建一个 MQD（Memory Queue Descriptor，队列描述符），并映射一个 doorbell（门铃）——你往队列里写完包，敲一下门铃，硬件的命令处理器（CP/MES）就来取包执行。

一句话：你说的"流"，本质是"一个能敲门铃的 AQL 用户态队列"。理解了这条，你就理解了 ROCm"用户态直接提交、绕开内核"的高性能设计精髓。

5、翻译链三：event 与 signal 的对应

• HIP 里你用 hipEvent_t 做计时、做流之间的等待。
• ROCr 把它实现成 hsa_signal_t——一个可以被"等待/唤醒"的信号值，AQL barrier 包也靠它表达依赖。
• libhsakmt 用 hsaKmtCreateEvent 创建一个 HsaEvent。
• KFD 在内核里管理这个事件，配合 GPU 中断在任务完成时唤醒等待方。

一句话：hipEvent → hsa_signal → HsaEvent → 内核事件+中断，是一条完整的"同步信号"翻译链。

6、进阶补充：其余概念对照

上面六类是主干。下面这些偏细，用得着再查——不影响主线阅读。

更多概念的四层对照（pinned / 拷贝 / 内存池 / 地址空间 / 多卡 / 代码对象）：

概念	HIP 层	ROCr / HSA 层	libhsakmt 层	KFD 内核
主机 pinned 内存	hipHostMalloc / hipHostRegister	hsa_amd_memory_lock	hsaKmtRegisterMemory	userptr
主机↔设备拷贝	hipMemcpy(Async)	hsa_amd_memory_async_copy	（提交到 SDMA 队列）	SDMA 引擎
内存池 / 区域	（隐式）	hsa_amd_memory_pool_t / region	HsaMemFlags、heap 类型	VRAM / GTT / USERPTR
地址空间 / 映射	裸指针	flat address	hsaKmtMapMemoryToGPU（GPUVA）	GPUVM、VA range
多卡 P2P 互访	hipDeviceEnablePeerAccess	hsa_amd_agents_allow_access	多 node MapMemoryToGPU	map 到多个 gpu
代码对象 / 模块	hipModule_t	hsa_executable_t / code object	—	—

几个值得留意的点：

• 拷贝其实也走队列：hipMemcpyAsync 底层是提交到 SDMA（DMA 拷贝引擎） 的队列，和 kernel 走的计算队列是两套，这样拷贝和计算能并行。
• 统一内存往下就是 SVM：核心表里 hipMallocManaged / hipMemAdvise 那条链，正好接上《AMD vs Intel vs NVIDIA：三家异构内存管理（SVM）》讲的 KFD SVM + HMM，可以对着读。
• 地址空间：hsaKmtMapMemoryToGPU 这一步才是真正把内存映射进 GPU 虚拟地址空间（GPUVA）的关键动作，很多"能分配但 GPU 访问崩"的问题都出在这一步没做对。

7、为什么要分这么多层？

读到这你可能会问：一个显存分配搞四层，是不是过度设计？其实每层都有它非留不可的理由：

1. HIP 层为了可移植：把 CUDA 生态的代码几乎零成本搬到 AMD，靠的就是 HIP 这层 API 兼容。
2. HSA 层为了标准化：agent/queue/signal 是开放的 HSA 标准，不绑死某家硬件，CPU、GPU、乃至别的加速器都能统一建模。
3. Thunk 层为了解耦：让上面的运行时不用直接跟内核 ioctl 打交道，内核接口变了，改这一薄层就行。
4. KFD 层为了"内核只做该做的事"：内核只提供最小必要的机制（建队列、分配内存、转发中断），策略都留在用户态——这正是 ROCm 高性能的根基（用户态直接提交命令，不用每次都陷入内核）。

一句话：这四层不是叠床架屋，而是"可移植 / 标准化 / 解耦 / 最小内核"四个目标各要一层。

8、想往下深挖：每层对应的专栏

这篇是"地图"，每一层的"逐行源码级实现"我在这些专栏里拆过：

想深入哪层	推荐专栏
ROCr + libhsakmt 源码	ROCm rocr 与 libhsakmt(thunk) 源码剖析
ROCm 整体机制	深探 ROCm
KFD 内存对象（BO）	AMD KFD BO 内存管理剖析
统一内存（SVM）	AMD KFD SVM 原理与实现剖析
驱动栈全景	AMDGPU ROCm 驱动栈全解析

9、写在最后

回到开头那两行代码。现在你应该能一眼看穿：hipMalloc 是 HIP → hsa_amd_memory_pool_allocate → hsaKmtAllocMemory → KFD 建 BO；hipLaunchKernel 是 HIP → 写 AQL 包 → 敲 doorbell → CP 取包执行。同一个东西每层换个名字，是因为每层站的抽象高度不一样——把这套"改名规律"记住，ROCm 源码读起来就从"迷宫"变成了"地图"。

如果这张对照表对你有用，欢迎收藏，下次读 ROCm 源码时对着查。

👉 觉得有用的话，收藏一下，转给同样在啃 ROCm 的同学。
👉 这几条翻译链里，你最想让我把哪一条拆到源码级？评论区告诉我。