前文：AMD rocr-libhsakmt分析系列6-1：共享机制-export分析了export hsa的bo，本文分析import graphics传过来的bo。
概述
函数名称: hsaKmtRegisterSharedHandle/hsakmt_fmm_register_graphics_handle
功能: 将图形栈（通过 DMA-BUF）导出的 GPU 内存注册到 ROCm/HSA 栈，实现图形与计算的内存互操作
1. 函数签名与参数
HSAKMT_STATUS hsakmt_fmm_register_graphics_handle(
HSAuint64 GraphicsResourceHandle, // DMA-BUF 文件描述符
HsaGraphicsResourceInfo *GraphicsResourceInfo, // 输出：资源信息
uint32_t *gpu_id_array, // GPU ID 数组（可选）
uint32_t gpu_id_array_size, // GPU ID 数组大小
HSA_REGISTER_MEM_FLAGS RegisterFlags // 注册标志
);参数详解
| 参数 | 类型 | 方向 | 说明 |
|---|---|---|---|
GraphicsResourceHandle | HSAuint64 | IN | 图形栈导出的 DMA-BUF fd |
GraphicsResourceInfo | HsaGraphicsResourceInfo* | OUT | 返回的资源信息（地址、大小、元数据） |
gpu_id_array | uint32_t* | IN | 要映射的 GPU ID 列表（可为 NULL） |
gpu_id_array_size | uint32_t | IN | GPU ID 数组大小 |
RegisterFlags | HSA_REGISTER_MEM_FLAGS | IN | 注册标志（如是否需要虚拟地址） |
2. 函数主流程分析
2.1 流程概览
┌──────────────────────────────────────────────────────────────────┐
│ 1. 参数验证 │
│ - 检查 gpu_id_array 与 gpu_id_array_size 的一致性 │
└────────────────┬─────────────────────────────────────────────────┘│▼
┌──────────────────────────────────────────────────────────────────┐
│ 2. 第一个 IOCTL: AMDKFD_IOC_GET_DMABUF_INFO │
│ - 输入: DMA-BUF fd │
│ - 输出: Buffer 信息（大小、GPU ID、标志、元数据） │
│ - 特性: 支持元数据动态调整（两次调用机制） │
└────────────────┬─────────────────────────────────────────────────┘│▼
┌──────────────────────────────────────────────────────────────────┐
│ 3. 选择内存 Aperture │
│ - 无 VA 要求 → mem_handle_aperture │
│ - SVM 需求 → svm.dgpu_aperture │
│ - 普通情况 → gpu_mem[id].gpuvm_aperture │
└────────────────┬─────────────────────────────────────────────────┘│▼
┌──────────────────────────────────────────────────────────────────┐
│ 4. 虚拟地址分配 │
│ - aperture_allocate_area_aligned() │
│ - 对齐要求: IMAGE_ALIGN (256KB) │
│ - 互斥锁保护: aperture->fmm_mutex │
└────────────────┬─────────────────────────────────────────────────┘│▼
┌──────────────────────────────────────────────────────────────────┐
│ 5. 第二个 IOCTL: AMDKFD_IOC_IMPORT_DMABUF │
│ - 输入: DMA-BUF fd, VA 地址, GPU ID │
│ - 输出: KFD handle │
│ - 功能: 将 DMA-BUF 导入到 KFD 并映射到指定 VA │
└────────────────┬─────────────────────────────────────────────────┘│▼
┌──────────────────────────────────────────────────────────────────┐
│ 6. 创建 VM 对象并注册 │
│ - aperture_allocate_object() │
│ - 设置内存标志（强制 CoarseGrain） │
│ - 保存元数据和 GPU ID 数组 │
└────────────────┬─────────────────────────────────────────────────┘│▼
┌──────────────────────────────────────────────────────────────────┐
│ 7. 填充输出信息 │
│ - MemoryAddress: 分配的虚拟地址 │
│ - SizeInBytes: Buffer 大小 │
│ - Metadata: 元数据指针 │
│ - NodeId: GPU 节点 ID │
└──────────────────────────────────────────────────────────────────┘3. 核心 IOCTL 分析
3.1 第一个 IOCTL: AMDKFD_IOC_GET_DMABUF_INFO
3.1.1 IOCTL 参数结构
struct kfd_ioctl_get_dmabuf_info_args {
__u64 size; // [OUT] Buffer 大小（字节）
__u64 metadata_ptr; // [IN/OUT] 元数据缓冲区指针
__u32 metadata_size; // [IN/OUT] 元数据大小
// IN: 用户分配的缓冲区大小
// OUT: 实际元数据大小
__u32 gpu_id; // [OUT] 创建此 buffer 的 GPU ID
__u32 flags; // [OUT] 内存标志 (KFD_IOC_ALLOC_MEM_FLAGS_*)
__u32 dmabuf_fd; // [IN] DMA-BUF 文件描述符
};3.1.2 实现代码
// 准备参数
infoArgs.dmabuf_fd = GraphicsResourceHandle;
infoArgs.metadata_size = GRAPHICS_METADATA_DEFAULT_SIZE; // 预设默认大小
metadata = calloc(infoArgs.metadata_size, 1);
infoArgs.metadata_ptr = (uint64_t)metadata;
r = hsakmt_ioctl(hsakmt_kfd_fd, AMDKFD_IOC_GET_DMABUF_INFO, (void *)&infoArgs);3.1.3 IOCTL 功能详解
作用:
- 验证 DMA-BUF: 检查 fd 是否有效且为 AMD GPU buffer
- 获取 Buffer 属性:
size: Buffer 的实际大小gpu_id: 创建该 buffer 的 GPU 标识flags: 内存属性标志（可写、一致性等）
- 提取元数据: 读取图形栈设置的元数据
3.1.4 内核侧处理（概念）
KFD 内核驱动侧:
1. 通过 dmabuf_fd 找到对应的 dma_buf 对象
2. 从 dma_buf->priv 获取 amdgpu_bo (buffer object)
3. 读取 BO 属性:- bo->tbo.base.size → size- bo->preferred_domains → flags- bo->metadata → 拷贝到用户空间
4. 确定所属 GPU ID
5. 返回给用户空间3.2 Aperture 选择逻辑
// 找到对应的 GPU 内存管理器
gpu_mem_id = gpu_mem_find_by_gpu_id(infoArgs.gpu_id);
// 根据需求选择 Aperture
if (!gpu_id_array && gpu_id_array_size == 0 &&
!RegisterFlags.ui32.requiresVAddr) {
// 场景 1: 不需要虚拟地址的导入
aperture = &mem_handle_aperture;
} else if (hsakmt_topology_is_svm_needed(
gpu_mem[gpu_mem_id].EngineId)) {
// 场景 2: 需要 SVM (Shared Virtual Memory)
aperture = svm.dgpu_aperture;
} else {
// 场景 3: 普通 GPU 虚拟内存
aperture = &gpu_mem[gpu_mem_id].gpuvm_aperture;
aperture_base = aperture->base;
}Aperture 类型说明
| Aperture 类型 | 使用场景 | 特点 |
|---|---|---|
mem_handle_aperture | 仅需要 handle，不需要 VA | 节省虚拟地址空间 |
svm.dgpu_aperture | SVM 应用（CPU/GPU 共享地址） | 地址对所有设备可见 |
gpuvm_aperture | 普通 GPU 专用内存 | 每个 GPU 独立地址空间 |
Aperture 选择影响:
- 性能: SVM 可能有额外的一致性开销
- 地址空间: 不同 aperture 管理不同的地址范围
- 可访问性: 决定哪些设备可以访问该内存
3.3 虚拟地址分配
// 在选定的 aperture 中分配虚拟地址
mem = aperture_allocate_area_aligned(
aperture, // 目标 aperture
NULL, // 地址提示（NULL = 自动选择）
infoArgs.size, // 大小
IMAGE_ALIGN // 对齐 (256KB)
);
}关键技术点:
- 线程安全: 使用互斥锁保护 aperture 操作
- 对齐要求: 256KB 对齐
- 符合 GPU 页表要求
- 优化图形资源访问性能
- 可能与纹理 tiling 要求相关
- 自动分配: 传递
NULL让系统自动选择最佳地址
内部机制 :
aperture_allocate_area_aligned() {
// 在 aperture 的空闲区域列表中查找
// 考虑对齐要求
// 返回分配的虚拟地址
}3.4 第二个 IOCTL: AMDKFD_IOC_IMPORT_DMABUF
3.4.1 IOCTL 参数结构
struct kfd_ioctl_import_dmabuf_args {
__u64 va_addr; // [IN] 目标虚拟地址
// 如果为 0，则不需要映射到 VA
__u64 handle; // [OUT] KFD 内部 handle
__u32 gpu_id; // [IN] 目标 GPU ID
__u32 dmabuf_fd; // [IN] DMA-BUF 文件描述符
};3.4.2 实现代码
// 准备导入参数
if (aperture == &mem_handle_aperture)
importArgs.va_addr = 0; // 不需要 VA 映射
else
importArgs.va_addr = VOID_PTRS_SUB(mem, aperture_base); // 计算相对地址
importArgs.gpu_id = infoArgs.gpu_id;
importArgs.dmabuf_fd = GraphicsResourceHandle;
// 调用 IOCTL
r = hsakmt_ioctl(hsakmt_kfd_fd, AMDKFD_IOC_IMPORT_DMABUF, (void *)&importArgs);
if (r) {
pthread_mutex_unlock(&aperture->fmm_mutex);
goto error_release_aperture;
}3.4.3 VA 地址计算
// VOID_PTRS_SUB 宏定义（概念）
#define VOID_PTRS_SUB(p1, p2) ((char*)(p1) - (char*)(p2))
// 示例:
// aperture_base = 0x8000000000
// mem = 0x8000100000
// va_addr = 0x8000100000 - 0x8000000000 = 0x100000 (相对偏移)为什么使用相对地址:
- GPU 页表中记录的是相对于 aperture base 的偏移
- 内核驱动会将相对地址转换为实际物理地址映射
- 支持 aperture 的动态重定位
3.4.4 IOCTL 功能详解
内核侧处理流程:
KFD 内核驱动侧:
┌─────────────────────────────────────────────────────────────┐
│ 1. DMA-BUF 查找 │
│ - dma_buf = dma_buf_get(dmabuf_fd) │
│ - amdgpu_bo = dma_buf->priv │
└────────────────┬────────────────────────────────────────────┘│▼
┌─────────────────────────────────────────────────────────────┐
│ 2. BO 导入到 KFD │
│ - 增加 BO 引用计数 │
│ - 创建 KFD 内部的 kfd_bo_va_list 结构 │
│ - 分配 handle (唯一标识符) │
└────────────────┬────────────────────────────────────────────┘│▼
┌─────────────────────────────────────────────────────────────┐
│ 3. GPU 页表映射 (如果 va_addr != 0) │
│ - 计算物理地址: bo->tbo.resource->start │
│ - 更新 GPU 页表: VA → PA 映射 │
│ - 设置页表项属性（缓存策略、权限等） │
└────────────────┬────────────────────────────────────────────┘│▼
┌─────────────────────────────────────────────────────────────┐
│ 4. 返回 handle │
│ - 用户空间通过 handle 管理该映射 │
│ - Unmap/Free 时需要提供此 handle │
└─────────────────────────────────────────────────────────────┘关键操作:
- 引用计数管理: 防止 BO 在使用时被释放
- 页表更新: 建立 GPU 可访问的虚拟地址映射
- Handle 生成: 提供用户空间管理接口
3.5 VM 对象创建与注册
// 转换内存标志
mflags = fmm_translate_ioc_to_hsa_flags(infoArgs.flags);
mflags.ui32.CoarseGrain = 1; // 强制设置为粗粒度内存
// 创建 VM 对象
obj = aperture_allocate_object(
aperture, // 所属 aperture
mem, // 虚拟地址
importArgs.handle, // KFD handle
infoArgs.size, // 大小
mflags // 内存标志
);关键点 - 强制 CoarseGrain:
mflags.ui32.CoarseGrain = 1; // ⚠️ 无条件设置为什么强制 CoarseGrain:
- 图形内存特性: 图形栈分配的内存通常是粗粒度的
- 性能考虑: 粗粒度内存带宽更高，适合图形资源
- 一致性简化: 避免细粒度一致性维护开销
3.6 填充输出信息
GraphicsResourceInfo->MemoryAddress = mem;
GraphicsResourceInfo->SizeInBytes = infoArgs.size;
GraphicsResourceInfo->Metadata = (void *)(unsigned long)infoArgs.metadata_ptr;
GraphicsResourceInfo->MetadataSizeInBytes = infoArgs.metadata_size;
hsakmt_gpuid_to_nodeid(infoArgs.gpu_id, &GraphicsResourceInfo->NodeId);
return HSAKMT_STATUS_SUCCESS;返回信息汇总:
| 字段 | 来源 | 说明 |
|---|---|---|
MemoryAddress | mem | 分配的虚拟地址（ROCm 栈可见） |
SizeInBytes | infoArgs.size | Buffer 大小 |
Metadata | infoArgs.metadata_ptr | 元数据指针 |
MetadataSizeInBytes | infoArgs.metadata_size | 元数据大小 |
NodeId | 从 gpu_id 转换 | HSA 节点 ID |
4. 关键技术点总结
4.1 两个 IOCTL 的协同
AMDKFD_IOC_GET_DMABUF_INFO AMDKFD_IOC_IMPORT_DMABUF│ │├─ 获取 buffer 属性 ├─ 导入 buffer 到 KFD├─ 验证 DMA-BUF 有效性 ├─ 建立 VA → PA 映射├─ 读取元数据 ├─ 更新 GPU 页表└─ 返回 GPU ID、大小、标志 └─ 返回 KFD handle│ │└────────── 信息传递 ──────────────┘(gpu_id, size, flags)为什么需要两个 IOCTL:
- 职责分离: 查询验证 vs 执行操作
- 灵活性: 用户可以先查询再决定是否导入
- 性能优化: 批量查询多个 buffer 后一次性导入
- 错误处理: 查询失败不影响系统状态
4.2 DMA-BUF 的生命周期管理
图形栈 用户空间 (Thunk) KFD 内核│ │ ││ amdgpu_bo_export │ ││─────────────────────────>│ ││ (返回 dmabuf_fd) │ ││ │ GET_DMABUF_INFO ││ │────────────────────────>││ │ │ dma_buf_get(fd)│ │ │ 引用计数 +1│ │<────────────────────────││ │ (返回 buffer 信息) ││ │ ││ │ IMPORT_DMABUF ││ │────────────────────────>││ │ │ dma_buf_get(fd)│ │ │ 引用计数 +1│ │<────────────────────────││ │ (返回 handle) ││ │ ││ close(dmabuf_fd) │ ││<─────────────────────────│ ││ │ │ 引用计数 -1│ │ │ (仍然 > 0)│ amdgpu_bo_free │ ││<─────────────────────────│ ││ │ │ 引用计数 -1│ │ │ (仍然 > 0)│ │ FREE_MEMORY_OF_GPU ││ │────────────────────────>││ │ │ 引用计数 -1│ │ │ (降为 0，释放)关键点:
- DMA-BUF fd 可以在导入后立即关闭
- KFD 持有独立的引用，保证 buffer 不会被过早释放
- 只有所有引用都释放后，物理内存才真正回收
4.3 内存标志的重要性
// 从 KFD 标志转换
mflags = fmm_translate_ioc_to_hsa_flags(infoArgs.flags);
// ⚠️ 强制设置 CoarseGrain
mflags.ui32.CoarseGrain = 1;内存标志影响的方面:
- 缓存策略: CoarseGrain vs FineGrain
- 一致性模型: 是否需要缓存刷新
- 访问权限: ReadOnly vs Writable
- 性能: 粗粒度内存带宽更高
5. 应用场景示例
场景 1: Vulkan 计算后处理
// Vulkan 渲染输出
VkDeviceMemory vkMem = RenderScene();
// 导出为 DMA-BUF
int fd = vkGetMemoryFdKHR(vkMem);
// 注册到 ROCm
HsaGraphicsResourceInfo info;
hsakmt_fmm_register_graphics_handle(
fd, &info, NULL, 0, {0}
);
// HIP kernel 进行后处理
PostProcessKernel<<<...>>>(info.MemoryAddress, ...);// 继续 Vulkan 渲染（无需拷贝）场景 2: 视频解码 + AI 分析
// VA-API 硬件解码
VASurfaceID surface = DecodeH264Frame();
// 导出帧
VADRMPRIMESurfaceDescriptor desc;
vaExportSurfaceHandle(surface, &desc);
// 注册到 ROCm
HsaGraphicsResourceInfo info;
hsakmt_fmm_register_graphics_handle(
desc.objects[0].fd, &info, NULL, 0, {0}
);
// 直接在 GPU 上运行推理
RunObjectDetection(model, info.MemoryAddress);6. 总结与改进建议
6.1 核心功能
hsakmt_fmm_register_graphics_handle 函数是 ROCm 图形互操作的关键入口点，通过两个 IOCTL 实现：
- GET_DMABUF_INFO: 查询 DMA-BUF buffer 的属性和元数据
- IMPORT_DMABUF: 将 buffer 导入到 KFD 并建立虚拟地址映射
6.2 改进建议
纯属在工作中遇到的问题，不一定合理。官方不支持，只能自己加了。
建议 1: 添加性能计数器
struct {
uint64_t total_imports;
uint64_t metadata_reallocs;
uint64_t avg_import_time_us;
} import_stats;建议 2: 支持批量导入
HSAKMT_STATUS hsakmt_fmm_register_graphics_handles_batch(
HSAuint64 *handles,
uint32_t count,
HsaGraphicsResourceInfo *infos
);建议 3: 异步导入支持
// 适用于大量小 buffer 的场景
HSAKMT_STATUS hsakmt_fmm_register_graphics_handle_async(
HSAuint64 handle,
completion_callback_t callback
);