rocm_bandwidth_test中的TransferBench (TB) 插件实现分析
ROCm Bandwidth Test 核心引擎 TransferBench 架构解析 TransferBench (TB) 1.60 是 ROCm Bandwidth Test (RBT) 的核心带宽测试引擎，采用插件化架构集成。系统由插件封装层和核心实现层组成，通过 plugin_main_entry 桥接原有可执行程序功能。核心执行流程包括环境配置收集、参数解析、测试运行和结果处理。关键 A
·
📋 概述
TransferBench (TB) 是 ROCm Bandwidth Test 的核心带宽测试引擎，版本 1.60。它被封装为插件形式集成到 RBT 框架中。
🏗️ 整体架构
rocm_bandwidth_test
├── plugins/tb/ # TB 插件封装层
│ ├── src/
│ │ ├── plugin_transferbench.cpp # 插件入口点
│ │ └── view_transferbench.cpp # 视图层（UI相关）
│ ├── include/
│ │ ├── plugin_transferbench.hpp
│ │ └── view_transferbench.hpp
│ └── transferbench/ # TB 客户端封装
│ ├── src/Client.cpp # 命令行解析和测试执行
│ └── include/
│ ├── Client.hpp
│ ├── Topology.hpp
│ ├── EnvVars.hpp
│ └── plugin_tb_link.hpp # 插件链接桥梁
└── deps/external/TransferBench/ # TB 核心实现
└── src/header/TransferBench.hpp # 单头文件实现
🔌 插件集成机制
1. 插件注册（plugin_transferbench.cpp）
// 宏注册插件元数据
AMD_WORK_BENCH_PLUGIN_SETUP(
"tb", // 插件名称
"Linux System Tools Team @AMD", // 作者
"Builtin: TransferBench", // 描述
"1.0.0" // 版本
)
{
using namespace wb_plugin_tb;
register_plugin_view(); // 注册视图组件
}
2. 插件入口点
// 插件主入口
int32_t plugin_main(int argc, char** argv)
{
// 调用 TransferBench 原始 main() 函数（重命名为 plugin_main_entry）
auto plugin_run_result = plugin_main_entry(argc, argv);
return plugin_run_result;
}
// TransferBench 原始主函数
extern int plugin_main_entry(int argc, char** argv);
设计思路：
- 将原有的独立可执行程序的
main()改名为plugin_main_entry() - 保持原有功能不变，便于从静态可执行文件转换为动态插件
🎯 核心执行流程
Client.cpp 主流程
int plugin_main_entry(int argc, char** argv)
{
1. 收集环境变量
├── EnvVars ev;
└── 读取各种配置选项
2. 处理命令行参数
├── if (argc <= 1) → 显示拓扑和帮助
├── 解析传输字节数 (numBytesPerTransfer)
└── 处理单位后缀 (K/M/G)
3. 运行预设测试（可选）
└── RunPreset(ev, numBytesPerTransfer, argc, argv)
4. 读取配置
├── 从配置文件读取
│ └── std::ifstream cfgFile(argv[1])
└── 或从命令行读取
└── if (!strcmp(argv[1], "cmdline"))
5. 主测试循环
for each line in config:
├── ParseTransfers(line, transfers) # 解析传输配置
├── 处理可变子执行器
└── for each byte size:
├── TransferBench::RunTransfers() # 执行传输
├── PrintResults() # 打印结果
└── PrintErrors() # 错误处理
}
🔑 核心 API 调用
1. TransferBench::ParseTransfers()
功能： 解析传输配置字符串
ErrResult ParseTransfers(
std::string str, // 配置字符串
std::vector<Transfer>& transfers // 输出：传输列表
);
配置格式示例：
# 语法: Src -> Executor:NumSubExecs -> Dst
G0 -> G0:8 -> G0 # GPU0 内部传输，8个子执行器
G0 -> C0:1 -> C0 # GPU0 -> CPU0
2. TransferBench::RunTransfers()
功能： 执行一组并行传输
bool RunTransfers(
ConfigOptions const& config, // 配置选项
vector<Transfer> const& transfers, // 传输列表
TestResults& results // 输出：测试结果
);
返回： true 表示成功，false 表示有致命错误
3. TransferBench::AllocateMemory()
功能： 分配指定类型的内存
static ErrResult AllocateMemory(
MemDevice memDevice, // 内存设备（类型+索引）
size_t numBytes, // 字节数
void** memPtr // 输出：内存指针
);
支持的内存类型：
MEM_CPU- 粗粒度固定 CPU 内存MEM_GPU- 粗粒度 GPU 全局内存MEM_CPU_FINE- 细粒度固定 CPU 内存MEM_GPU_FINE- 细粒度 GPU 全局内存MEM_CPU_UNPINNED- 未固定 CPU 内存MEM_MANAGED- 统一内存
4. TransferBench::GetNumExecutors()
功能： 获取可用执行器数量
int GetNumExecutors(ExeType exeType);
执行器类型：
EXE_CPU© - CPU 执行器（子执行器 = CPU 线程）EXE_GPU_GFX(G) - GPU 内核执行器（子执行器 = CU/线程块）EXE_GPU_DMA(D) - GPU SDMA 执行器EXE_NIC(I) - NIC RDMA 执行器EXE_NIC_NEAREST(N) - 最近的 NIC RDMA 执行器
5. TransferBench::GetNumSubExecutors()
功能： 获取子执行器数量
int GetNumSubExecutors(ExeDevice exeDevice);
示例：
{EXE_GPU_GFX, 0}→ 返回 GPU 0 的 CU 数量（如 40）{EXE_CPU, 0}→ 返回 CPU 线程数
📊 关键数据结构
Transfer（传输定义）
struct Transfer
{
size_t numBytes; // 传输字节数
vector<MemDevice> srcs; // 源内存设备列表
vector<MemDevice> dsts; // 目标内存设备列表
ExeDevice exeDevice; // 执行器
int32_t exeSubIndex; // 执行器子索引（如 XCD）
int numSubExecs; // 子执行器数量（如 CU 数量）
};
TestResults（测试结果）
struct TestResults
{
int numTimedIterations; // 计时迭代次数
size_t totalBytesTransferred; // 总传输字节数
double avgTotalDurationMsec; // 平均总持续时间（毫秒）
double avgTotalBandwidthGbPerSec; // 平均总带宽（GB/s）
double overheadMsec; // 开销时间
map<ExeDevice, ExeResult> exeResults; // 每个执行器的结果
vector<TransferResult> tfrResults; // 每个传输的结果
vector<ErrResult> errResults; // 错误/警告列表
};
ExeDevice（执行器设备）
struct ExeDevice
{
ExeType exeType; // 执行器类型（C/G/D/I/N）
int32_t exeIndex; // 执行器索引（设备编号）
};
MemDevice（内存设备）
struct MemDevice
{
MemType memType; // 内存类型（C/G/B/F/U/M）
int32_t memIndex; // 设备索引
};
⚙️ 配置选项
ConfigOptions（完整配置）
struct ConfigOptions
{
GeneralOptions general; // 通用选项
DataOptions data; // 数据选项
GfxOptions gfx; // GPU 图形执行器选项
DmaOptions dma; // DMA 执行器选项
NicOptions nic; // NIC 执行器选项
};
GeneralOptions（通用选项）
struct GeneralOptions
{
int numIterations = 10; // 计时迭代次数（负数表示按秒运行）
int numSubIterations = 1; // 每次迭代的子迭代次数
int numWarmups = 3; // 预热迭代次数
int recordPerIteration = 0; // 记录每次迭代的计时信息
int useInteractive = 0; // 在传输循环前暂停等待用户输入
};
GfxOptions（GPU 执行器选项）
struct GfxOptions
{
int blockSize = 256; // 线程块大小（必须是 64 的倍数）
vector<uint32_t> cuMask = {}; // CU 掩码位向量
vector<vector<int>> prefXccTable = {}; // [src][dst] 首选 XCD 表
int unrollFactor = 4; // 内核展开因子
int useHipEvents = 1; // 使用 HIP 事件计时
int useMultiStream = 0; // 使用多流
int useSingleTeam = 0; // 所有子执行器组成单个团队
int waveOrder = 0; // 波前排序
};
🔄 执行流程详解
1. 配置解析阶段
配置字符串: "G0 -> G0:8 -> G0 1048576"
↓
ParseTransfers()
↓
Transfer {
numBytes = 1048576, // 1 MB
srcs = [{MEM_GPU, 0}], // GPU 0 内存
dsts = [{MEM_GPU, 0}], // GPU 0 内存
exeDevice = {EXE_GPU_GFX, 0}, // GPU 0 执行器
numSubExecs = 8 // 8 个 CU
}
2. 内存分配阶段
RunTransfers() 内部：
for each Transfer:
for each src in Transfer.srcs:
AllocateMemory(src, numBytes, &srcMem)
for each dst in Transfer.dsts:
AllocateMemory(dst, numBytes, &dstMem)
3. 传输执行阶段
循环流程：
Warmup iterations (不计时):
for i = 0 to numWarmups:
ExecuteTransfer()
Timed iterations (计时):
StartTimer()
for i = 0 to numIterations:
ExecuteTransfer()
if recordPerIteration:
RecordIterationTime()
StopTimer()
4. 结果计算阶段
计算带宽：
avgDurationMsec = totalTime / numIterations
avgBandwidthGbPerSec = (numBytes / 1.0E9) / avgDurationMsec * 1000.0
🎨 执行器类型详解
GPU GFX 执行器（EXE_GPU_GFX）
特点：
- 使用 GPU 内核（kernel）执行传输
- 子执行器 = CU（计算单元）或线程块
- 支持 CU 掩码精细控制
- 支持多 XCD（多芯片设计）
内核参数：
- blockSize: 线程块大小（256）
- unrollFactor: 循环展开因子（4）
- waveOrder: 波前调度顺序
- cuMask: 位掩码指定使用哪些 CU
GPU DMA 执行器（EXE_GPU_DMA）
特点：
- 使用 SDMA 引擎执行传输
- 不涉及计算单元，纯硬件 DMA
- 支持 HIP 或 HSA 拷贝
选项：
- useHipEvents: 使用 HIP 事件计时
- useHsaCopy: 使用 HSA 拷贝（而非 HIP）
CPU 执行器（EXE_CPU）
特点：
- 使用 CPU 线程执行 memcpy
- 子执行器 = CPU 线程
- 支持 NUMA 绑定
📈 结果输出格式
标准输出格式
Test 1:
Executor: GFX 00 | 230.456 GB/s | 4.532 ms | 1048576 bytes | 230.456 GB/s (sum)
Transfer 00 | 230.456 GB/s | 4.532 ms | 1048576 bytes | G0 -> G00:8 -> G0
Aggregate (CPU) | 230.456 GB/s | 4.532 ms | 1048576 bytes | Overhead: 0.123 ms
CSV 输出格式
Test#,Transfer#,NumBytes,Src,Exe,Dst,CUs,BW(GB/s),Time(ms),SrcAddr,DstAddr
1,0,1048576,G0,G0,G0,8,230.456,4.532,0x7f1234,0x7f5678
🛠️ 重要辅助函数
1. GetClosestCpuNumaToGpu()
int GetClosestCpuNumaToGpu(int gpuIndex);
返回离 GPU 最近的 NUMA 节点索引。
2. GetClosestNicToGpu()
int GetClosestNicToGpu(int gpuIndex);
返回离 GPU 最近的 NIC 索引（用于 RDMA）。
3. GetNumExecutorSubIndices()
int GetNumExecutorSubIndices(ExeDevice exeDevice);
返回执行器子索引数量（如 XCD 数量、DMA 引擎数量）。
🔍 环境变量控制
TransferBench 通过 EnvVars 类读取环境变量：
struct EnvVars
{
ConfigOptions ToConfigOptions(); // 转换为配置选项
void DisplayEnvVars(); // 显示环境变量
static void DisplayUsage(); // 显示使用说明
// 成员变量（部分）
int outputToCsv; // CSV 输出模式
int showIterations; // 显示每次迭代
int samplingFactor; // 采样因子
int minNumVarSubExec; // 最小可变子执行器数
int maxNumVarSubExec; // 最大可变子执行器数
};
🧪 测试流程示例
示例 1: GPU 内部带宽测试
# 配置
G0 -> G0:40 -> G0 1048576
# 执行流程
1. ParseTransfers() 解析配置
2. AllocateMemory({MEM_GPU, 0}, 1MB, &src)
3. AllocateMemory({MEM_GPU, 0}, 1MB, &dst)
4. 预热 3 次迭代
5. 计时执行 10 次迭代
6. 计算平均带宽
7. 打印结果: ~230 GB/s
8. DeallocateMemory()
示例 2: Host-to-Device 带宽测试
# 配置
C0 -> G0:8 -> G0 1048576
# 执行流程
1. AllocateMemory({MEM_CPU, 0}, 1MB, &src) # 固定 CPU 内存
2. AllocateMemory({MEM_GPU, 0}, 1MB, &dst) # GPU 内存
3. GPU 0 的 8 个 CU 执行传输
4. 结果: ~25 GB/s (PCIe 4.0 x16 理论值)
🚨 错误处理
ErrResult 结构
enum ErrType {
ERR_NONE = 0, // 无错误
ERR_WARN = 1, // 警告（结果可能不准确）
ERR_FATAL = 2 // 致命错误（结果无效）
};
struct ErrResult {
ErrType errType;
std::string errMsg;
};
错误检查宏
#define ERR_CHECK(expression) \
do { \
ErrResult err = (expression); \
if (err.errType == ERR_FATAL) return err; \
} while(0)
使用示例：
ERR_CHECK(AllocateMemory(device, size, &ptr));
ERR_CHECK(RunTransfers(config, transfers, results));
📚 典型使用场景
1. 基准测试
./rocm_bandwidth_test test_simple.cfg
2. P2P 测试（如果支持）
./rocm_bandwidth_test test_p2p.cfg
3. 数据大小扫描
./rocm_bandwidth_test test.cfg 0 # 0 表示扫描多种大小
4. 预设测试
./rocm_bandwidth_test <preset_name> [bytes]
🎯 设计亮点
1. 单头文件实现
- 整个 TransferBench 核心在单个
.hpp文件中 - 便于集成和分发
- 降低编译依赖
2. 插件化架构
- 原本独立的工具转换为插件
- 保持向后兼容
- 易于集成到更大的测试框架
3. 灵活的执行器模型
- 支持多种执行器类型（CPU/GPU/DMA/NIC）
- 子执行器概念支持精细控制
- 可变子执行器数量自动优化
4. 全面的内存类型支持
- 粗粒度/细粒度
- 固定/未固定
- 统一内存
- 支持各种组合测试
5. 详细的性能分析
- 每次迭代计时
- 每个执行器带宽
- 每个传输带宽
- 聚合带宽和开销分析
🔧 关键实现技术
HIP/ROCm API
// 设备管理
hipGetDeviceCount(&numGpus);
hipSetDevice(gpuIndex);
// 内存管理
hipMalloc(&ptr, size);
hipHostMalloc(&ptr, size, hipHostMallocDefault);
hipMemcpy(dst, src, size, hipMemcpyDeviceToDevice);
// P2P 访问
hipDeviceCanAccessPeer(&canAccess, gpu0, gpu1);
hipDeviceEnablePeerAccess(gpu1, 0);
// 计时
hipEventCreate(&start);
hipEventRecord(start);
hipEventElapsedTime(&ms, start, stop);
HSA API
// 用于更底层的控制
hsa_agent_t agent;
hsa_amd_memory_pool_t pool;
hsa_amd_memory_copy(dst, src, size);
NUMA 库
// NUMA 节点管理
numa_available();
numa_bind(node);
numa_alloc_onnode(size, node);
📝 总结
TransferBench 是一个功能强大、设计精良的带宽测试工具：
核心优势：
- ✅ 支持多种执行器和内存类型
- ✅ 灵活的配置系统
- ✅ 详细的性能分析
- ✅ 插件化集成
- ✅ 全面的错误处理
关键 API：
ParseTransfers()- 配置解析RunTransfers()- 执行传输AllocateMemory()- 内存分配GetNumExecutors()- 查询执行器
适用场景：
- GPU 内部带宽测试
- Host-Device 传输测试
- P2P 带宽测试（专业卡）
- 多 GPU 拓扑分析
- 性能基准测试
有“AI”的1024 = 2048，欢迎大家加入2048 AI社区
更多推荐
6
0- 0
已为社区贡献3条内容
所有评论(0)