06 - rocrtst 性能测试详解

本文档介绍rocrtst性能测试套件（suites/performance/）中的各个模块，帮助你理解测试指标、运行方法和结果解读。

1. 性能测试概览

性能测试注册在rocrtstPerf测试套件下，共 6 个源码模块。运行所有性能测试：

sudorocrtst64--gtest_filter="rocrtstPerf.*"

2. dispatch_time — 分发延迟

源文件：dispatch_time.h/dispatch_time.cc

测量从 CPU 提交空 kernel 到 GPU 执行完成的端到端延迟。

2.1 测试用例

2.2 运行

# 运行所有分发延迟测试sudorocrtst64--gtest_filter="rocrtstPerf.AQL_Dispatch_Time_*"# 仅运行单 kernel 自旋等待（最低延迟基线）sudorocrtst64--gtest_filter="rocrtstPerf.AQL_Dispatch_Time_Single_SpinWait"

2.3 测量原理

CPU 时间线： ├─ T1: 写入 AQL 包到队列 ────┐ │ │ ← Dispatch Latency ├─ T2: Signal 完成通知 ─────┘ │ Dispatch Time = T2 - T1

• SpinWait：CPU 通过hsa_signal_wait_scacquire()自旋轮询 Signal 值变化，延迟最低但消耗 CPU
• Interrupt：CPU 通过中断方式等待，延迟略高但 CPU 友好

2.4 结果解读

输出示例：

Mean dispatch latency (single, spin wait): 5.23 us Mean dispatch latency (single, interrupt): 8.45 us

涉及 Kernel：dispatch_time_kernels.cl（空 kernel）

3. enqueueLatency — 入队延迟

源文件：enqueueLatency.h/enqueueLatency.cc

测量将 AQL 数据包写入队列的 CPU 端延迟（不包含 GPU 执行时间）。

3.1 测试用例

该测试内部创建两个实例：

• EnqueueLatency(true)— 单包入队
• EnqueueLatency(false)— 多包入队

3.2 运行

sudorocrtst64--gtest_filter="rocrtstPerf.ENQUEUE_LATENCY"

3.3 测量原理

CPU 时间线： ├─ T1: 准备 AQL 包 ────────┐ │ │ ← Enqueue Latency ├─ T2: 写入队列门铃寄存器 ───┘ │ Enqueue Time = T2 - T1

与dispatch_time的区别：

• enqueueLatency仅测量 CPU 端写包的耗时
• dispatch_time测量从写包到 GPU 执行完成的全程

3.4 结果解读

4. memory_async_copy — 异步内存拷贝

源文件：memory_async_copy.h/memory_async_copy.cc

测量不同路径下的异步内存拷贝带宽，是最重要的性能测试之一。

4.1 测试用例

4.2 拷贝路径类型

4.3 运行

# 运行默认拷贝测试（CPU↔GPU + GPU↔GPU）sudorocrtst64--gtest_filter="rocrtstPerf.Memory_Async_Copy"# 增加迭代次数以获得更稳定的数据sudorocrtst64--gtest_filter="rocrtstPerf.Memory_Async_Copy"-i20# 开启详细输出查看中间数据sudorocrtst64--gtest_filter="rocrtstPerf.Memory_Async_Copy"-v2

4.4 结果解读

输出示例：

Copy: CPU(0) -> GPU(1) Size: 4096 KB Time: 0.52 ms BW: 7.69 GB/s Copy: GPU(1) -> CPU(0) Size: 4096 KB Time: 0.48 ms BW: 8.33 GB/s Copy: GPU(1) -> GPU(2) Size: 4096 KB Time: 0.21 ms BW: 19.05 GB/s

涉及 API：hsa_amd_memory_async_copy()

5. memory_async_copy_on_engine — 指定引擎拷贝

源文件：memory_async_copy_on_engine.h/memory_async_copy_on_engine.cc

继承自MemoryAsyncCopy，在指定的 DMA 引擎上执行拷贝并验证数据正确性。

5.1 测试用例

5.2 运行

sudorocrtst64--gtest_filter="rocrtstPerf.Memory_Async_Copy_On_Engine"

5.3 说明

• 使用hsa_amd_memory_async_copy_on_engine()API（AMD 扩展）
• 可以指定使用哪个 SDMA（System DMA）引擎
• 包含数据验证（Benchmark with Verification），确保拷贝正确性

6. memory_async_copy_numa — NUMA 感知拷贝

源文件：memory_async_copy_numa.h/memory_async_copy_numa.cc

测试 NUMA 架构下，不同 NUMA 节点间的内存拷贝带宽。

6.1 测试用例

⚠️ 此测试默认 DISABLED，且构建时通过ENABLE_COPY_NUMA=OFF排除源文件。启用需要修改 CMake 选项。

6.2 启用方法

cmake-DENABLE_COPY_NUMA=ON...

运行时：

sudorocrtst64--gtest_also_run_disabled_tests--gtest_filter="*NUMA*"

6.3 依赖

• 需要libhwloc（NUMA 拓扑检测）
• 多 NUMA 节点系统（如双路服务器 + 多 GPU）

7. agent_preload — Agent 预加载

源文件：agent_preload.h/agent_preload.cc

测量 Agent 预加载对性能的影响，对比开启/关闭预加载时的操作延迟。

7.1 测试用例

7.2 运行

sudorocrtst64--gtest_filter="rocrtstPerf.Agent_Preload_Latency"

7.3 测量内容

该测试对比两个场景的延迟：

预加载（Preload）会在 Agent 初始化时预先加载 Blit kernel，避免首次拷贝时的冷启动开销。

7.4 结果解读

输出示例：

Profiling Enable Latency: Without preload: 1500.00 us With preload: 200.00 us Improvement: 1300.00 us First Async Copy Latency: Without preload: 2000.00 us With preload: 300.00 us Improvement: 1700.00 us

8. 性能测试最佳实践

8.1 获取稳定结果

# 增加迭代次数sudorocrtst64--gtest_filter="rocrtstPerf.*"-i50# 多次重复整个测试sudorocrtst64--gtest_filter="rocrtstPerf.*"-i20--gtest_repeat=3

8.2 减少干扰因素

# 固定 CPU 频率（避免频率调节影响）sudocpupower frequency-set-gperformance# 固定 GPU 频率sudorocm-smi--setperfdeterminism1800# 关闭不必要的进程

8.3 结合 GPU 监控

# 打印每个测试前后的 GPU 状态（温度、频率、功耗等）sudorocrtst64--gtest_filter="rocrtstPerf.*"-m1

8.4 导出结果

# 导出为 XML（可被 CI 系统解析）sudorocrtst64--gtest_filter="rocrtstPerf.*"-i20--gtest_output=xml:perf_report.xml

9. 性能测试速查表

上一篇：05-功能测试详解
下一篇：07-负面测试与压力测试详解 — 异常场景与高并发测试