KFD 驱动如何管理 Compute Queue（计算队列）

在 AMDGPU + ROCm 架构中，KFD（Kernel Fusion Driver，内核融合驱动） 是 amdgpu 内核驱动的一部分，专门用于支持异构计算（Heterogeneous Compute），

其核心任务之一就是管理用户态提交的 Compute Queue（计算队列）。Compute Queue 是 GPU 执行计算内核（如 HIP 或 OpenCL kernel）的命令通道。KFD 负责

在内核中为每个队列分配资源、建立与硬件的连接，并确保安全、高效地调度到 GPU 上执行。

 

一、关键术语说明

　　KFD：Kernel Fusion Driver（内核融合驱动）

　　Compute Queue：计算队列，用于提交计算任务到 GPU

　　BO：Buffer Object（缓冲区对象），DRM/GPU 内存管理的基本单位

　　MQD：Memory-mapped Queue Descriptor（内存映射队列描述符），GPU 硬件用于描述一个队列状态的数据结构

　　GEM：Graphics Execution Manager（图形执行管理器），DRM 子系统中的内存对象抽象层

　　PASID：Process Address Space ID（进程地址空间标识符），用于 SVM（共享虚拟内存）场景下的进程隔离

　　VMID：Virtual Memory ID（虚拟内存标识符），GPU MMU 使用的地址空间标签

　　SVM：Shared Virtual Memory（共享虚拟内存），CPU 与 GPU 共享同一虚拟地址空间

　　ACE：Asynchronous Compute Engine（异步计算引擎），AMD RDNA 架构中的硬件调度单元

 

二、Compute Queue 的创建流程（含 BO 分配）

当用户空间的 HSA Runtime（如 ROCr）调用 hsa_queue_create() 时，会通过 ioctl 向 /dev/kfd 发起 KFD_IOC_CREATE_QUEUE 请求。KFD 内核驱动处理该请求，主要步骤如下：

1. 分配 Ring Buffer（环形缓冲区）

　　Ring Buffer 用于存放 AQL（Architected Queuing Language）命令包。

　　KFD 调用 amdgpu_gem_new() 创建一个 Buffer Object（BO），作为 ring buffer 的底层存储。

　　　　这个 BO 可位于 GTT（Graphics Translation Table）（即系统内存，GPU 可访问）或 VRAM（显存），取决于策略和硬件能力。

　　　　该 BO 被映射到进程的 SVM 地址空间，用户态可直接读写。

 

BO（Buffer Object）是什么？

BO 是 DRM 子系统中表示一块 GPU 可访问内存的抽象对象。它封装了物理内存、GPU 虚拟地址、引用计数、缓存属性等信息。

在 amdgpu 中，BO 由 TTM（Translation Table Maps）或 GEM 后端管理。

 

2. 分配 Doorbell（门铃）资源

　　Doorbell 是一小块特殊内存区域（通常 8 字节），用户写入后会触发 GPU 硬件中断或唤醒命令处理器。

　　KFD 从 GPU 的 doorbell aperture（门铃地址窗口） 中分配一个 slot。

　　该 doorbell 通过 mmap(/dev/kfd) 映射到用户空间，用户可直接写入更新 WPTR（Write Pointer）

3. 分配 MQD（Memory-mapped Queue Descriptor）的 BO

　　MQD 是 GPU 硬件定义的数据结构，描述了一个队列的所有状态（如 ring 基地址、大小、读写指针、CU 掩码等）。

　　KFD 为 MQD 单独分配一个 Buffer Object（BO），通常较小（几百字节），存放于 GTT 或 VRAM。

　　调用 ASIC 特定的 MQD 管理器（如 mqd_v10.c 对应 GFX10/RDNA2）的 init_mqd() 函数，填充 MQD 内容。

为什么 MQD 也要用 BO？

因为 MQD 必须位于 GPU 可访问的内存中，且其地址需写入 GPU 寄存器（如 CP_MQD_BASE_ADDR）。使用 BO 可统一管理生命周期、地址映射和内存类型。

 

4. 注册到 QPD（Queue Process Device）

　　每个进程在每个 GPU 上有一个 QPD（Queue Process Device） 结构。

　　新创建的队列被加入 QPD 的队列列表。

　　若是该进程在此 GPU 上的第一个队列，KFD 还会：

　　　　分配 PASID（Process Address Space ID）　　

　　　　分配 VMID（Virtual Memory ID）

　　　　将进程的页表基地址（PTBASE）绑定到 VMID，实现 SVM

5. 激活队列（Load to Hardware）

　　队列默认懒激活（lazy activation）。

　　当用户首次写 doorbell 提交命令时，KFD 调用 mqd_manager->load_mqd()：

　　　　通过 MMIO（Memory-Mapped I/O） 写寄存器，或

　　　　通过 间接加载机制（如 ACE 的 MQD list）

　　　　将 MQD 的 GPU 虚拟地址告知硬件调度器（如 ACE）

　　GPU 的 Compute Command Processor 开始监控该 ring buffer 的 WPTR 变化，并执行命令。

 

三、硬件交互（以 RDNA2/GFX10 为例）

　　MQD 格式由 AMD 公开部分规范，包含 ring_base_addr、ring_size、wptr、cu_mask 等字段。

　　Doorbell 写入 → 触发 GPU 的 Doorbell 引擎 → 更新内部 WPTR → 唤醒 CP。

　　ACE（Asynchronous Compute Engine） 硬件自动轮询多个活跃队列，实现高并发调度。

 

四、总结

KFD 通过以下方式管理 Compute Queue：

功能	实现方式
内存分配	使用 Buffer Object（BO） 表示 ring buffer 和 MQD
地址空间	基于 PASID + VMID 实现 SVM（共享虚拟内存）
硬件对接	通过 MQD（Memory-mapped Queue Descriptor） 描述队列状态
低延迟提交	用户态直接写 doorbell，无需 syscall
多队列调度	由 GPU 硬件 ACE 自动轮询，KFD 负责注册/注销

 

这种设计使得 ROCm 应用能够以极低开销将计算任务提交到 AMD GPU，充分发挥其大规模并行计算能力。