kfdtest: KFDGraphicsInterop测试用例功能分析-优快云博客

文档概述

文件路径: KFDGraphicsInterop.cpp
测试目的: 验证 ROCm 计算栈与图形渲染栈之间的内存互操作性

1. 背景知识

1.1 为什么需要图形互操作

在 AMD GPU 上，存在两个独立的软件栈：

图形渲染栈: Mesa/libdrm → amdgpu kernel driver
计算/ROCm 栈: HIP/HSA → KFD kernel driver

许多应用场景需要两个栈共享 GPU 内存：

机器学习可视化: 计算结果直接渲染显示
视频处理: 解码后的帧用于计算分析
科学计算: 计算结果导出到图形应用
游戏 AI: 推理结果影响渲染场景

1.3 DMA-BUF 机制

DMA-BUF (DMA Buffer Sharing) 是 Linux 内核提供的跨子系统内存共享框架：

图形栈分配内存 → 导出为 DMA-BUF fd → 计算栈导入 fd → 访问同一块物理内存

2. 测试用例详解

2.1 测试用例 1: RegisterGraphicsHandle

2.1.1 测试目标

验证通过图形栈（libdrm/amdgpu）分配的 GPU 内存可以被 ROCm/HSA 栈正确注册、映射和访问。

2.1.2 测试流程图

┌─────────────────────────────────────────────────────────────────┐
│ 1. 环境检查                                                       │
│    - 检查 GPU 型号（跳过 Tonga）                                   │
│    - 查找 DRM render node                                        │
└────────────────┬────────────────────────────────────────────────┘
                 │
                 ▼
┌─────────────────────────────────────────────────────────────────┐
│ 2. 通过图形栈分配 VRAM                                            │
│    amdgpu_bo_alloc_request:                                     │
│    - alloc_size: PAGE_SIZE (4KB) 或 8*PAGE_SIZE                 │
│    - preferred_heap: AMDGPU_GEM_DOMAIN_VRAM                     │
│    - flags: AMDGPU_GEM_CREATE_CPU_ACCESS_REQUIRED               │
└────────────────┬────────────────────────────────────────────────┘
                 │
                 ▼
┌─────────────────────────────────────────────────────────────────┐
│ 3. CPU 访问并设置元数据                                            │
│    - amdgpu_bo_cpu_map() 映射到 CPU 地址空间                      │
│    - memset(pCpuMap, 0xaa, PAGE_SIZE) 填充测试数据                │
│    - amdgpu_bo_set_metadata() 设置元数据                          │
│    - amdgpu_bo_export() 导出为 DMA-BUF fd                        │
└────────────────┬────────────────────────────────────────────────┘
                 │
                 ▼
┌─────────────────────────────────────────────────────────────────┐
│ 4. 注册到 ROCm/HSA 栈                                            │
│    hsaKmtRegisterGraphicsHandleToNodes(dmabufFd, &info, ...)    │
│    - 输入: DMA-BUF 文件描述符                                     │
│    - 输出: HsaGraphicsResourceInfo (地址、大小、元数据)             │
│    - KFD 驱动增加 BO 引用计数                                      │
└────────────────┬────────────────────────────────────────────────┘
                 │
                 ▼
┌─────────────────────────────────────────────────────────────────┐
│ 5. 验证注册信息                                                   │
│    - info.SizeInBytes == alloc.alloc_size ✓                     │
│    - info.MetadataSizeInBytes == metadata_size ✓                │
│    - info.Metadata == 原始元数据 ✓                               │
└────────────────┬────────────────────────────────────────────────┘
                 │
                 ▼
┌─────────────────────────────────────────────────────────────────┐
│ 6. GPU 访问测试                                                  │
│    - hsaKmtMapMemoryToGPU() 映射到 GPU 地址空间                   │
│    - 编译并加载 CopyDwordIsa kernel                              │
│    - dispatch.SetArgs(图形内存地址, 目标buffer)                   │
│    - 提交到 PM4 队列执行                                          │
│    - 同步等待完成                                                 │
└────────────────┬────────────────────────────────────────────────┘
                 │
                 ▼
┌─────────────────────────────────────────────────────────────────┐
│ 7. 验证拷贝结果                                                   │
│    EXPECT_EQ(dstBuffer[0], 0xaaaaaaaa) ✓                        │
│    - 确认 GPU 成功读取了图形栈分配的内存                             │
└────────────────┬────────────────────────────────────────────────┘
                 │
                 ▼
┌─────────────────────────────────────────────────────────────────┐
│ 8. 查询指针信息                                                   │
│    hsaKmtQueryPointerInfo(info.MemoryAddress, &ptrInfo)         │
│    验证:                                                         │
│    - ptrInfo.Type == HSA_POINTER_REGISTERED_GRAPHICS ✓          │
│    - ptrInfo.MemFlags.ui32.CoarseGrain == 1 ✓                   │
│    - ptrInfo.Node == 当前 GPU 节点 ✓                             │
└────────────────┬────────────────────────────────────────────────┘
                 │
                 ▼
┌─────────────────────────────────────────────────────────────────┐
│ 9. 清理资源                                                      │
│    - hsaKmtUnmapMemoryToGPU() 解除 GPU 映射                      │
│    - hsaKmtDeregisterMemory() 注销内存                           │
│    - close(dmabufFd) 关闭文件描述符                               │
│    - amdgpu_bo_free(handle) 释放 BO                             │
└─────────────────────────────────────────────────────────────────┘

2.1.3 关键代码段分析

步骤 1: 内存分配

struct amdgpu_bo_alloc_request alloc;
alloc.alloc_size = size;
alloc.preferred_heap = AMDGPU_GEM_DOMAIN_VRAM;  // 分配在 VRAM
alloc.flags = AMDGPU_GEM_CREATE_CPU_ACCESS_REQUIRED;  // 需要 CPU 可访问

步骤 3: 元数据设置

const char metadata[] = "This data is really meta.";
struct amdgpu_bo_metadata meta;
meta.size_metadata = strlen(metadata) + 1;
memcpy(meta.umd_metadata, metadata, metadata_size);
amdgpu_bo_set_metadata(handle, &meta);

用途: 测试元数据在栈之间的传递
实际应用: 可存储纹理格式、对齐要求等信息

步骤 4: DMA-BUF 导出与注册

uint32_t dmabufFd;
amdgpu_bo_export(handle, amdgpu_bo_handle_type_dma_buf_fd, &dmabufFd);

HsaGraphicsResourceInfo info;
hsaKmtRegisterGraphicsHandleToNodes(dmabufFd, &info, 1, nodes);

close(dmabufFd);      // KFD 已持有引用，可以关闭 fd
amdgpu_bo_free(handle);  // KFD 已持有引用，可以释放 BO

关键机制: KFD 在注册时会增加 BO 的内核引用计数
资源管理: 原始 fd 和 handle 可以立即释放

步骤 6: GPU Kernel 执行

Assembler* m_pAsm = pKFDGraphicsInterop->GetAssemblerFromNodeId(gpuNode);
m_pAsm->RunAssembleBuf(CopyDwordIsa, isaBuffer.As<char*>());

Dispatch dispatch(isaBuffer);
dispatch.SetArgs(info.MemoryAddress, dstBuffer.As<void*>());
dispatch.Submit(queue);
dispatch.Sync(g_TestTimeOut);

CopyDwordIsa: 简单的 DWORD 拷贝 shader
验证目标: 确认 GPU 能读取图形栈分配的内存

步骤 8: 指针类型验证

EXPECT_EQ(ptrInfo.Type, HSA_POINTER_REGISTERED_GRAPHICS);
EXPECT_EQ(ptrInfo.MemFlags.ui32.CoarseGrain, 1);

关键发现: 注册的图形内存被标记为 粗粒度 (Coarse Grain)
性能影响: 粗粒度内存不支持细粒度原子操作，但带宽更高

2.1.4 测试覆盖的功能点

功能点	验证方法	预期结果
DMA-BUF 导入	`hsaKmtRegisterGraphicsHandleToNodes`	成功返回资源信息
元数据传递	比较 `info.Metadata` 和原始元数据	完全一致
大小信息	比较 `info.SizeInBytes` 和分配大小	完全一致
GPU 内存映射	`hsaKmtMapMemoryToGPU`	映射成功
GPU 读取访问	Kernel 拷贝并验证数据	数据正确 (0xaaaaaaaa)
指针类型	`hsaKmtQueryPointerInfo`	`HSA_POINTER_REGISTERED_GRAPHICS`
内存粒度	查询 `CoarseGrain` 标志	等于 1
资源清理	Unmap 和 Deregister	无内存泄漏

2. 关键技术要点

2.1 DMA-BUF 生命周期管理

图形栈                   内核                     ROCm栈
  |                       |                        |
  | amdgpu_bo_alloc       |                        |
  |---------------------->| BO refcount = 1        |
  |                       |                        |
  | amdgpu_bo_export      |                        |
  |---------------------->| 创建 DMA-BUF            |
  |<----------------------| BO refcount = 2        |
  | (返回 fd)              |                        |
  |                       |                        |
  |                       |   RegisterGraphics     |
  |                       |<-----------------------|
  |                       | BO refcount = 3        |
  |                       |----------------------->|
  |                       |   (返回地址)            |
  |                       |                        |
  | close(fd)             |                        |
  |---------------------->| BO refcount = 2        |
  |                       |                        |
  | amdgpu_bo_free        |                        |
  |---------------------->| BO refcount = 1        |
  |                       |                        |
  |                       |   DeregisterMemory     |
  |                       |<-----------------------|
  |                       | BO refcount = 0        |
  |                       | 释放物理内存             |

关键点:

KFD 在注册时会持有 BO 的内核引用
即使用户空间释放了所有句柄，BO 仍然有效
只有 Deregister 后才真正释放

3.2 内存类型和性能

内存类型	CoarseGrain	FineGrain	Uncached	带宽	延迟	用途
VRAM (Normal)	✅	❌	❌	最高	低	GPU 计算主力
VRAM (Fine-grained)	❌	✅	❌	高	中	需要原子操作
VRAM (Uncached)	❌	❌	✅	很低	高	特殊场景
注册图形内存	✅	❌	❌	高	低	图形/计算互操作
System Memory	❌	✅	❌	中	高	CPU/GPU 共享

测试验证的重点:

EXPECT_EQ(ptrInfo.MemFlags.ui32.CoarseGrain, 1);

确认注册的图形内存是 粗粒度 (CoarseGrain)
这是高带宽访问的关键

3. 测试执行与调试

3.1 运行测试

# 运行特定测试
./kfdtest --gtest_filter=KFDGraphicsInterop.RegisterGraphicsHandle

4.2 预期输出

成功输出:

[==========] Running 1 test from 1 test suite.
[----------] Global test environment set-up.
[----------] 1 test from KFDGraphicsInterop
[ RUN      ] KFDGraphicsInterop.RegisterGraphicsHandle
[       OK ] KFDGraphicsInterop.RegisterGraphicsHandle (XXX ms)
[----------] 1 test from KFDGraphicsInterop (XXX ms total)

[==========] 1 test from 1 test suite ran. (XXX ms total)
[  PASSED  ] 1 test.

失败场景:

找不到 render node: GPU 没有配置图形驱动
分配失败: VRAM 不足或不支持 CPU 访问
注册失败: KFD 驱动问题
数据不匹配: GPU 访问失败或数据损坏
CoarseGrain != 1: 内存类型错误

4.3 调试技巧

启用详细日志:

检查 DMA-BUF:

# 查看进程的 DMA-BUF
ls -la /proc/$(pidof kfdtest)/fd | grep dmabuf

# 查看 DMA-BUF 信息
cat /sys/kernel/debug/dma_buf/bufinfo

验证 GPU 内存:

# 检查 GPU VRAM 使用
rocm-smi --showmeminfo vram

# 检查 large BAR 支持
rocminfo | grep -A 10 "Marketing Name"

5. 实际应用场景

5.1 Vulkan + HIP 混合应用

场景: 游戏引擎中的光线追踪降噪

// Vulkan 渲染光线追踪原始输出
VkDeviceMemory vkImage = RenderRayTracing();

// 导出为 DMA-BUF
int dmabufFd = vkGetMemoryFdKHR(vkImage);

// 导入到 HIP/ROCm
HsaGraphicsResourceInfo info;
hsaKmtRegisterGraphicsHandleToNodes(dmabufFd, &info, ...);

// HIP kernel 处理降噪
hipLaunchKernel(DenoiseKernel, ..., info.MemoryAddress, ...);

// 结果直接用于后续渲染，无需拷贝

5.2 OpenGL + OpenCL 科学可视化

场景: 流体模拟实时可视化

// OpenGL 分配纹理
GLuint texture;
glGenTextures(1, &texture);
glBindTexture(GL_TEXTURE_2D, texture);

// 导出为 DMA-BUF
int dmabufFd = eglExportDMABUFImageMESA(texture);

// OpenCL 导入
cl_mem clImage = clCreateFromDMABUF(dmabufFd);

// OpenCL kernel 更新模拟
clEnqueueNDRangeKernel(queue, simulationKernel, ..., clImage, ...);

// OpenGL 直接渲染更新后的纹理，无需 CPU 拷贝
glDrawArrays(...);

5.3 视频编解码 + AI 推理

场景: 实时视频内容分析

// VA-API 硬件解码视频帧到 GPU 内存
VASurfaceID surface;
vaCreateSurfaces(..., &surface);
vaDecodePicture(..., surface);

// 导出解码后的帧
VADRMPRIMESurfaceDescriptor desc;
vaExportSurfaceHandle(surface, &desc);

// ROCm 导入用于推理
hsaKmtRegisterGraphicsHandleToNodes(desc.objects[0].fd, &info, ...);

// 直接在 GPU 上运行目标检测/分类
RunInference(model, info.MemoryAddress);

6. 总结

6.1 测试价值

RegisterGraphicsHandle 测试验证了 ROCm 栈的关键功能：

✅ DMA-BUF 跨子系统内存共享
✅ 图形栈分配的内存可被计算栈访问
✅ 元数据正确传递
✅ GPU 能正确读取注册的图形内存
✅ 内存类型标记正确（CoarseGrain）

附录 A: 相关 API 参考

A.1 libdrm (amdgpu) APIs

amdgpu_bo_alloc() - 分配 GPU buffer object
amdgpu_bo_cpu_map() - CPU 映射 BO
amdgpu_bo_set_metadata() - 设置 BO 元数据
amdgpu_bo_export() - 导出为 DMA-BUF fd
amdgpu_bo_free() - 释放 BO

A.2 libhsakmt APIs

hsaKmtRegisterGraphicsHandleToNodes() - 注册图形句柄
hsaKmtMapMemoryToGPU() - 映射到 GPU 地址空间
hsaKmtQueryPointerInfo() - 查询指针信息
hsaKmtUnmapMemoryToGPU() - 解除 GPU 映射
hsaKmtDeregisterMemory() - 注销内存

A.3 相关常量

// 内存域
AMDGPU_GEM_DOMAIN_VRAM  // GPU 本地内存
AMDGPU_GEM_DOMAIN_GTT   // 系统内存

// 分配标志
AMDGPU_GEM_CREATE_CPU_ACCESS_REQUIRED  // 需要 CPU 访问
AMDGPU_GEM_CREATE_NO_CPU_ACCESS        // 不需要 CPU 访问

// 指针类型
HSA_POINTER_REGISTERED_GRAPHICS  // 注册的图形内存
HSA_POINTER_ALLOCATED            // HSA 分配的内存

// 内存标志
CoarseGrain = 1  // 粗粒度 (高带宽)
FineGrain = 1    // 细粒度 (支持原子)
Uncached = 1     // 不缓存 (低带宽)