本文档介绍rocrtst性能测试套件suites/performance/中的各个模块帮助你理解测试指标、运行方法和结果解读。1. 性能测试概览性能测试注册在rocrtstPerf测试套件下共 6 个源码模块。运行所有性能测试sudorocrtst64--gtest_filterrocrtstPerf.*模块测量目标关键指标dispatch_timeGPU 任务分发延迟微秒 (μs)enqueueLatencyAQL 包入队延迟微秒 (μs)memory_async_copy异步内存拷贝带宽GB/smemory_async_copy_on_engine指定引擎的异步拷贝GB/smemory_async_copy_numaNUMA 感知的异步拷贝GB/sagent_preloadAgent 预加载延迟优化微秒 (μs)2. dispatch_time — 分发延迟源文件dispatch_time.h/dispatch_time.cc测量从 CPU 提交空 kernel 到 GPU 执行完成的端到端延迟。2.1 测试用例测试用例等待方式Kernel 数量说明AQL_Dispatch_Time_Single_SpinWait自旋等待单个最低延迟场景AQL_Dispatch_Time_Single_Interrupt中断等待单个中断驱动场景AQL_Dispatch_Time_Multi_SpinWait自旋等待多个批量分发自旋AQL_Dispatch_Time_Multi_Interrupt中断等待多个批量分发中断2.2 运行# 运行所有分发延迟测试sudorocrtst64--gtest_filterrocrtstPerf.AQL_Dispatch_Time_*# 仅运行单 kernel 自旋等待最低延迟基线sudorocrtst64--gtest_filterrocrtstPerf.AQL_Dispatch_Time_Single_SpinWait2.3 测量原理CPU 时间线 ├─ T1: 写入 AQL 包到队列 ────┐ │ │ ← Dispatch Latency ├─ T2: Signal 完成通知 ─────┘ │ Dispatch Time T2 - T1SpinWaitCPU 通过hsa_signal_wait_scacquire()自旋轮询 Signal 值变化延迟最低但消耗 CPUInterruptCPU 通过中断方式等待延迟略高但 CPU 友好2.4 结果解读输出示例Mean dispatch latency (single, spin wait): 5.23 us Mean dispatch latency (single, interrupt): 8.45 us指标典型值 (MI300X)说明Single SpinWait3-8 μs最优分发延迟基线Single Interrupt6-15 μs中断开销约 3-7 μsMulti SpinWait更低平均值流水线效应涉及 Kerneldispatch_time_kernels.cl空 kernel3. enqueueLatency — 入队延迟源文件enqueueLatency.h/enqueueLatency.cc测量将 AQL 数据包写入队列的 CPU 端延迟不包含 GPU 执行时间。3.1 测试用例测试用例说明ENQUEUE_LATENCY分别测量单包入队和多包入队延迟该测试内部创建两个实例EnqueueLatency(true)— 单包入队EnqueueLatency(false)— 多包入队3.2 运行sudorocrtst64--gtest_filterrocrtstPerf.ENQUEUE_LATENCY3.3 测量原理CPU 时间线 ├─ T1: 准备 AQL 包 ────────┐ │ │ ← Enqueue Latency ├─ T2: 写入队列门铃寄存器 ───┘ │ Enqueue Time T2 - T1与dispatch_time的区别enqueueLatency仅测量 CPU 端写包的耗时dispatch_time测量从写包到 GPU 执行完成的全程3.4 结果解读指标典型值说明Single packet 1 μs单个 AQL 包写入耗时Multi packet略高多包连续写入的平均耗时4. memory_async_copy — 异步内存拷贝源文件memory_async_copy.h/memory_async_copy.cc测量不同路径下的异步内存拷贝带宽是最重要的性能测试之一。4.1 测试用例测试用例说明Memory_Async_Copy测试多种拷贝路径的带宽4.2 拷贝路径类型路径类型缩写说明Host → DeviceH2D主机内存 → GPU 显存Device → HostD2HGPU 显存 → 主机内存Peer-to-PeerP2PGPU → GPU直连Host → Device (Remote)H2DRemote远端 NUMA 节点 → GPUDevice → Host (Remote)D2HRemoteGPU → 远端 NUMA 节点P2P (Remote)P2PRemoteGPU → GPU跨 NUMA4.3 运行# 运行默认拷贝测试CPU↔GPU GPU↔GPUsudorocrtst64--gtest_filterrocrtstPerf.Memory_Async_Copy# 增加迭代次数以获得更稳定的数据sudorocrtst64--gtest_filterrocrtstPerf.Memory_Async_Copy-i20# 开启详细输出查看中间数据sudorocrtst64--gtest_filterrocrtstPerf.Memory_Async_Copy-v24.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路径典型带宽取决于H2D / D2H10-26 GB/sPCIe Gen4/Gen5 带宽P2P直连50-200 GB/sxGMI / Infinity Fabric 带宽P2P跨 NUMA较低跨 socket 互联带宽涉及 APIhsa_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 测试用例测试用例说明Memory_Async_Copy_On_Engine在指定 SDMA 引擎上拷贝并验证5.2 运行sudorocrtst64--gtest_filterrocrtstPerf.Memory_Async_Copy_On_Engine5.3 说明使用hsa_amd_memory_async_copy_on_engine()APIAMD 扩展可以指定使用哪个 SDMASystem 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_Memory_Async_Copy_NUMANUMA 感知异步拷贝默认禁用⚠️ 此测试默认 DISABLED且构建时通过ENABLE_COPY_NUMAOFF排除源文件。启用需要修改 CMake 选项。6.2 启用方法cmake-DENABLE_COPY_NUMAON...运行时sudorocrtst64--gtest_also_run_disabled_tests--gtest_filter*NUMA*6.3 依赖需要libhwlocNUMA 拓扑检测多 NUMA 节点系统如双路服务器 多 GPU7. agent_preload — Agent 预加载源文件agent_preload.h/agent_preload.cc测量 Agent 预加载对性能的影响对比开启/关闭预加载时的操作延迟。7.1 测试用例测试用例说明Agent_Preload_Latency对比预加载开启/关闭的延迟7.2 运行sudorocrtst64--gtest_filterrocrtstPerf.Agent_Preload_Latency7.3 测量内容该测试对比两个场景的延迟场景说明Profiling Enable 延迟首次调用hsa_amd_profiling_async_copy_enable()的耗时首次异步拷贝延迟首次hsa_amd_memory_async_copy()的耗时包含 Blit 初始化预加载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 us8. 性能测试最佳实践8.1 获取稳定结果# 增加迭代次数sudorocrtst64--gtest_filterrocrtstPerf.*-i50# 多次重复整个测试sudorocrtst64--gtest_filterrocrtstPerf.*-i20--gtest_repeat38.2 减少干扰因素# 固定 CPU 频率避免频率调节影响sudocpupower frequency-set-gperformance# 固定 GPU 频率sudorocm-smi--setperfdeterminism1800# 关闭不必要的进程8.3 结合 GPU 监控# 打印每个测试前后的 GPU 状态温度、频率、功耗等sudorocrtst64--gtest_filterrocrtstPerf.*-m18.4 导出结果# 导出为 XML可被 CI 系统解析sudorocrtst64--gtest_filterrocrtstPerf.*-i20--gtest_outputxml:perf_report.xml9. 性能测试速查表命令说明--gtest_filterrocrtstPerf.*所有性能测试--gtest_filterrocrtstPerf.AQL_Dispatch_Time_*分发延迟--gtest_filterrocrtstPerf.ENQUEUE_LATENCY入队延迟--gtest_filterrocrtstPerf.Memory_Async_Copy异步拷贝带宽--gtest_filterrocrtstPerf.Memory_Async_Copy_On_Engine指定引擎拷贝--gtest_filterrocrtstPerf.Agent_Preload_Latency预加载延迟上一篇05-功能测试详解下一篇07-负面测试与压力测试详解 — 异常场景与高并发测试