ROCm带宽测试原理与优化
前文介绍了rocm_bandwidth_test工具功能和安装方法，本文详细分析带宽测试的原理。
1. 整体架构
ROCm Bandwidth Test (RBT-NG) 是下一代 ROCm 带宽测试工具，采用插件化架构，核心基于 TransferBench (TB) 引擎。
设计理念
- 模块化设计：通过插件系统实现功能扩展
- 多执行器支持：支持 CPU、GPU GFX、GPU DMA、NIC 等多种执行单元
- 灵活的内存类型：支持多种内存类型和设备组合
- 精确测量：提供纳秒级时间精度和详细的性能指标
2. 核心测试原理
2.1 Transfer（传输）概念
每个 Transfer 定义了一次数据传输操作，是测试的基本单元。
Transfer 数据结构：
struct Transfer {
size_t numBytes; // 传输的字节数
vector<MemDevice> srcs; // 源内存设备列表
vector<MemDevice> dsts; // 目标内存设备列表
ExeDevice exeDevice; // 执行器设备
int32_t exeSubIndex; // 执行器子索引（XCC/SDMA引擎）
int numSubExecs; // 子执行器数量
};
核心算法：
将所有源数据求和后写入所有目标位置
sum = src[0] + src[1] + ... + src[N-1]
dst[0] = dst[1] = ... = dst[M-1] = sum
特殊情况：
- 0 个源：对目标执行 memset 操作（填充固定值）
- 1 个源：执行简单的内存复制
- 多个源：执行累加操作后写入目标
2. 执行器类型 (Executor Types)
2.1 CPU 执行器 (EXE_CPU)
特点：
- 使用多线程并行执行
- 每个子执行器在独立的 CPU 线程上运行
- 基于 NUMA 节点进行内存分配和线程绑定
核心实现：
static void CpuReduceKernel(SubExecParam const& p, int numSubIterations) {
if (numSrcs == 0) {
// Memset 操作
memset(dst, MEMSET_CHAR, N * sizeof(float));
} else if (numSrcs == 1) {
// 单源复制
memcpy(dst, src, N * sizeof(float));
} else {
// 多源累加
for (int j = 0; j < N; j++) {
sum = src[0][j];
for (int i = 1; i < numSrcs; i++)
sum += src[i][j];
for (int i = 0; i < numDsts; i++)
dst[i][j] = sum;
}
}
}
执行流程：
- 绑定到指定的 NUMA 节点
- 为每个传输创建异步线程
- 每个传输内部启动多个子执行器线程
- 使用
std::thread实现并行
2.2 GPU GFX 执行器 (EXE_GPU_GFX)
特点：
- 使用 GPU 计算内核（HIP/CUDA kernel）执行
- 支持模板化设计，可配置 block size 和 unroll 因子
- 支持多种 wavefront 排序模式
- 支持 XCC（芯片组）指定和 CU 掩码
核心内核：
template <int BLOCKSIZE, int UNROLL>
__global__ void GpuReduceKernel(SubExecParam* params,
int waveOrder,
int numSubIterations) {
// 1. 记录开始时间戳
if (threadIdx.x == 0)
startCycle = GetTimestamp();
// 2. 过滤指定的 XCC
if (preferredXccId != -1 && xccId != preferredXccId)
return;
// 3. 向量化读取和处理（使用 float4）
float4 val[UNROLL];
for (int u = 0; u < UNROLL; u++)
val[u] = srcFloat4[0][idx + u * unrlStride * warpSize];
// 4. 累加多个源
for (int s = 1; s < numSrcs; s++)
for (int u = 0; u < UNROLL; u++)
val[u] += srcFloat4[s][idx + u * unrlStride * warpSize];
// 5. 写入所有目标
for (int d = 0; d < numDsts; d++)
for (int u = 0; u < UNROLL; u++)
dstFloat4[d][idx + u * unrlStride * warpSize] = val[u];
// 6. 记录结束时间戳
if (threadIdx.x == 0) {
p.stopCycle = GetTimestamp();
p.startCycle = startCycle;
}
}
关键优化技术：
-
向量化访问：使用
float4类型，一次读写 16 字节 -
循环展开：支持 1-8 的展开因子，减少循环开销
-
Wavefront 排序：6 种排序模式优化内存访问
U,W,C: Unroll, Wave, CUU,C,W: Unroll, CU, WaveW,U,C: Wave, Unroll, CUW,C,U: Wave, CU, UnrollC,U,W: CU, Unroll, WaveC,W,U: CU, Wave, Unroll
-
多流支持：可以为每个传输创建独立的 HIP 流
-
XCC 亲和性：在多芯片 GPU 上指定执行的芯片组
内核启动配置：
dim3 gridSize(xccDim, numSubExecs, 1);
dim3 blockSize(cfg.gfx.blockSize, 1);
// 支持 64, 128, 192, 256, 320, 384, 448, 512 的 block size
// 支持 1-8 的 unroll 因子
GpuKernelTable[blockSize/64 - 1][unrollFactor - 1]
<<<gridSize, blockSize, 0, stream>>>
(params, waveOrder, numSubIterations);
2.3 GPU DMA 执行器 (EXE_GPU_DMA)
特点：
- 使用 GPU 的 SDMA（System DMA）引擎
- 不占用 GPU 计算资源
- 支持指定特定的 SDMA 引擎
两种实现方式：
- HIP Copy（默认）：
hipMemcpyAsync(dst, src, numBytes, hipMemcpyDefault, stream);
- HSA Copy（高级模式）：
hsa_amd_memory_async_copy_on_engine(
dst, dstAgent,
src, srcAgent,
numBytes, 0, NULL,
signal, sdmaEngineId, true
);
SDMA 引擎选择：
- 可以通过
exeSubIndex指定特定的 SDMA 引擎 - 支持 SDMA_ID_0, SDMA_ID_1 等
- 引擎可用性通过
hsa_amd_memory_copy_engine_status查询
2.4 NIC 执行器 (EXE_NIC)
特点：
- 使用 InfiniBand Verbs 进行 RDMA 传输
- 支持 RoCE (RDMA over Converged Ethernet)
- 使用队列对（Queue Pairs）实现并行传输
核心实现：
// 1. 创建队列对
CreateQueuePair(cfg, protectDomain, completionQueue, queuePair);
// 2. 初始化队列对
InitQueuePair(queuePair, port, rdmaAccessFlags);
// 3. 转换到 RTR (Ready to Receive)
TransitionQpToRtr(queuePair, ...);
// 4. 转换到 RTS (Ready to Send)
TransitionQpToRts(queuePair);
// 5. 执行 RDMA Write
ibv_post_send(queuePair, &sendWorkRequest, &badWorkReq);
// 6. 轮询完成队列
while (numComplete < qpCount) {
nc = ibv_poll_cq(completionQueue, 1, &wc);
if (nc > 0) numComplete++;
}
支持特性：
- RoCE v1 和 v2
- IPv4 和 IPv6
- Relaxed Ordering 优化
- 多队列对并行传输
3. 内存类型 (Memory Types)
支持 7 种内存类型，适应不同的测试场景：
| 类型 | 代码 | 描述 | 特点 |
|---|---|---|---|
| MEM_CPU | C | 粗粒度固定 CPU 内存 | 高带宽，GPU 可访问 |
| MEM_GPU | G | 粗粒度 GPU 全局内存 | GPU 原生内存 |
| MEM_CPU_FINE | B | 细粒度固定 CPU 内存 | 缓存一致性 |
| MEM_GPU_FINE | F | 细粒度 GPU 全局内存 | 原子操作友好 |
| MEM_CPU_UNPINNED | U | 非固定 CPU 内存 | 可分页内存 |
| MEM_MANAGED | M | 统一管理内存 | 自动迁移 |
| MEM_NULL | N | 空内存 | 只读测试 |
内存分配实现：
static ErrResult AllocateMemory(MemDevice memDevice,
size_t numBytes,
void** memPtr) {
if (IsCpuMemType(memType)) {
// 设置 NUMA 策略
numa_set_preferred(memDevice.memIndex);
if (memType == MEM_CPU_FINE) {
hipHostMalloc(memPtr, numBytes, hipHostMallocNumaUser);
} else if (memType == MEM_CPU) {
hipHostMalloc(memPtr, numBytes,
hipHostMallocNumaUser | hipHostMallocNonCoherent);
} else if (memType == MEM_CPU_UNPINNED) {
*memPtr = numa_alloc_onnode(numBytes, memDevice.memIndex);
}
// 验证页面分配在正确的 NUMA 节点
CheckPages((char*)*memPtr, numBytes, memDevice.memIndex);
} else if (IsGpuMemType(memType)) {
hipSetDevice(memDevice.memIndex);
if (memType == MEM_GPU) {
hipMalloc(memPtr, numBytes);
} else if (memType == MEM_GPU_FINE) {
hipExtMallocWithFlags(memPtr, numBytes, hipDeviceMallocUncached);
} else if (memType == MEM_MANAGED) {
hipMallocManaged(memPtr, numBytes);
}
hipMemset(*memPtr, 0, numBytes);
}
}
3. 测试执行流程
阶段 1：准备阶段 (Preparation)
1.1 内存分配
for each Transfer:
for each source:
AllocateMemory(srcMemDevice, numBytes)
EnablePeerAccess(if needed)
for each destination:
AllocateMemory(dstMemDevice, numBytes)
EnablePeerAccess(if needed)
1.2 数据准备
// 生成参考数据
PrepareReference(cfg, srcBuffer, srcBufferIdx);
// 填充公式（伪随机但可重现）
element[i] = ((i * 517) % 383 + 31) * (srcBufferIdx + 1)
// 计算预期输出
for (numSrcs sources):
expectedOutput = sum(all sources)
1.3 子执行器配置
PrepareSubExecParams(cfg, transfer, resources);
// 数据分割策略
- 每个子执行器处理 blockBytes 的倍数
- 尽量均匀分配
- 最后一个子执行器处理剩余数据
// Single Team 模式
if (useSingleTeam):
所有子执行器协同处理整个数组（交错访问）
else:
每个子执行器处理独立的子数组
阶段 2：执行阶段 (Execution)
2.1 预热迭代
for i = 0 to numWarmups - 1:
RunAllTransfers(iteration = -1) // 不计时
预热缓存和 GPU 内核
2.2 计时迭代
for iteration = 0 to numIterations - 1:
startTime = now()
// 并行执行所有 Executor
for each Executor:
async_execute:
for each Transfer:
if (useMultiStream):
LaunchInSeparateStream()
else:
CombineInSingleKernel()
WaitForAllExecutors()
endTime = now()
totalDuration += (endTime - startTime)
if (alwaysValidate):
ValidateResults()
2.3 时间测量机制
CPU 执行器：
auto start = std::chrono::high_resolution_clock::now();
ExecuteTransfer();
auto end = std::chrono::high_resolution_clock::now();
duration = std::chrono::duration_cast<std::chrono::duration<double>>(end - start);
GPU GFX 执行器（HIP Events）：
hipEventRecord(startEvent, stream);
LaunchKernel<<<...>>>(params);
hipEventRecord(stopEvent, stream);
hipEventElapsedTime(&duration, startEvent, stopEvent);
GPU GFX 执行器（设备时间戳）：
// 在内核内部
__global__ void Kernel() {
if (threadIdx.x == 0) {
startCycle = wall_clock64();
}
// 执行数据传输
...
__syncthreads();
if (threadIdx.x == 0) {
stopCycle = wall_clock64();
}
}
// 在主机端
duration_ms = (stopCycle - startCycle) / wallClockRate;
阶段 3：验证阶段 (Validation)
ValidateAllTransfers(cfg, transfers, resources, reference, outputBuffer);
for each Transfer:
for each destination:
if (validateDirect or IsCpuMemType):
output = dstMemPtr
else:
hipMemcpy(outputBuffer, dstMemPtr, numBytes, hipMemcpyDefault)
output = outputBuffer
if (memcmp(output, expected, numBytes) != 0):
// 找到第一个不匹配的位置
for i in range(N):
if (output[i] != expected[i]):
ReportError(transferIdx, i, expected[i], output[i])
return FATAL_ERROR
4. 带宽计算方法
4.1 单次传输带宽
TransferBandwidth (GB/s) = (numBytes / 1.0e6) / avgDurationMsec
示例：
- 传输 1 GB 数据
- 耗时 10 ms
- 带宽 = (1,000,000,000 / 1,000,000) / 10 = 100 GB/s
4.2 执行器带宽
ExecutorBandwidth (GB/s) = totalBytes / avgDurationMsec
其中：
totalBytes = sum(所有该执行器处理的传输字节数)
avgDurationMsec = totalDuration / numIterations
4.3 系统总带宽
SystemBandwidth (GB/s) = totalBytesTransferred / avgTotalDurationMsec
其中：
totalBytesTransferred = sum(所有传输的字节数)
avgTotalDurationMsec = 墙上时钟时间 / numIterations
4.4 开销计算
Overhead (ms) = avgTotalDuration - min(所有执行器的持续时间)
开销表示系统同步、调度等非数据传输的时间消耗。
4.5 每迭代统计（可选）
当启用 recordPerIteration 时：
for each iteration:
记录：
- perIterMsec[i]: 该迭代的持续时间
- perIterCUs[i]: 使用的计算单元 (XCC:CU)
计算：
- 标准差：variance from avgBandwidth
- 最小/最大带宽
- 带宽稳定性
5. 性能优化技术
5.1 数据分块 (Data Blocking)
原理：
blockBytes = 256; // 默认值，可配置
每个子执行器处理的数据量为 blockBytes 的倍数
优势：
- 缓存行对齐（通常 64 或 128 字节）
- 减少跨缓存行访问
- 提高内存访问效率
5.2 向量化访问
GPU 实现：
// 使用 float4 代替 float
float4* srcFloat4 = (float4*)src;
float4* dstFloat4 = (float4*)dst;
float4 val = srcFloat4[idx]; // 一次读取 16 字节
dstFloat4[idx] = val; // 一次写入 16 字节
性能提升：
- 减少 4 倍的内存事务数量
- 提高内存总线利用率
- 现代 GPU 对向量访问优化良好
5.3 循环展开 (Loop Unrolling)
实现：
template <int UNROLL>
__global__ void Kernel() {
float4 val[UNROLL]; // 寄存器数组
#pragma unroll
for (int u = 0; u < UNROLL; u++) {
val[u] = src[idx + u * stride];
}
#pragma unroll
for (int u = 0; u < UNROLL; u++) {
dst[idx + u * stride] = val[u];
}
}
优势：
- 减少循环控制开销
- 提高指令级并行 (ILP)
- 隐藏内存延迟
- 支持 1-8 的展开因子
5.4 Wavefront 排序优化
6 种排序模式：
// Mode 0: U,W,C - 适合连续访问
unroll, wave, cu
stride = 1, UNROLL, UNROLL * nWaves
// Mode 1: U,C,W - 适合跨 CU 分布
unroll, cu, wave
stride = 1, UNROLL, UNROLL * nTeams
// Mode 2: W,U,C - Wave 优先
wave, unroll, cu
stride = 1, nWaves, nWaves * UNROLL
// Mode 3: W,C,U - Wave 和 CU 优先
wave, cu, unroll
stride = 1, nWaves, nWaves * nTeams
// Mode 4: C,U,W - CU 优先
cu, unroll, wave
stride = 1, nTeams, nTeams * UNROLL
// Mode 5: C,W,U - CU 和 Wave 优先
cu, wave, unroll
stride = 1, nTeams, nTeams * nWaves
选择策略：
- 连续访问：模式 0 或 2
- 分散访问：模式 1 或 4
- 负载均衡：模式 3 或 5
5.5 流并行 (Stream Parallelism)
单流模式（默认）：
// 所有传输合并到一个内核启动
dim3 gridSize(xccDim, totalSubExecs, 1);
Kernel<<<gridSize, blockSize, 0, stream>>>(allParams);
多流模式：
// 每个传输使用独立的流
for each Transfer:
Kernel<<<gridSize, blockSize, 0, stream[i]>>>(params[i]);
优势：
- 提高 GPU 利用率
- 允许并发执行
- 更好的硬件队列利用
- 需要设置
GPU_MAX_HW_QUEUES环境变量
5.6 P2P 和 Large BAR 优化
P2P 启用：
EnablePeerAccess(int deviceId, int peerDeviceId) {
hipDeviceCanAccessPeer(&canAccess, deviceId, peerDeviceId);
if (canAccess) {
hipDeviceEnablePeerAccess(peerDeviceId, 0);
}
}
Large BAR 检查：
for each GPU:
hipDeviceGetAttribute(&isLargeBar,
hipDeviceAttributeIsLargeBar,
gpuIndex);
if (!isLargeBar) {
ReportWarning("Large BAR not enabled for GPU %d", gpuIndex);
}
优势：
- 直接 GPU-to-GPU 传输
- 无需 CPU 中转
- 降低延迟，提高带宽
5.7 NUMA 优化
CPU 线程绑定：
numa_run_on_node(numaNode); // 在指定 NUMA 节点运行
内存分配策略：
numa_set_preferred(numaNode); // 设置首选 NUMA 节点
ptr = numa_alloc_onnode(size, node); // 在指定节点分配
页面验证：
CheckPages(ptr, numBytes, targetNumaNode) {
move_pages(0, numPages, pages, NULL, status, 0);
for each page:
if (status[i] != targetNumaNode) {
ReportError("Page not on correct NUMA node");
}
}
6. 测试场景支持
6.1 单向传输模式
Host-to-Device (H2D)
Transfer: C0 -> G0 (CPU内存 -> GPU内存)
Executor: CPU, GPU_GFX, 或 GPU_DMA
示例：
1 1 C0 G0 G0 // 使用 GPU_GFX
1 1 C0 D0 G0 // 使用 GPU_DMA
Device-to-Host (D2H)
Transfer: G0 -> C0 (GPU内存 -> CPU内存)
Executor: CPU, GPU_GFX, 或 GPU_DMA
示例：
1 1 G0 G0 C0 // 使用 GPU_GFX
1 1 G0 D0 C0 // 使用 GPU_DMA
Device-to-Device (D2D)
Transfer: G0 -> G1 (GPU0 -> GPU1)
需要：P2P 支持，Large BAR
示例：
1 1 G0 G0 G1 // GPU0 执行
1 1 G0 G1 G1 // GPU1 执行
RDMA Transfer
Transfer: G0 -> RemoteGPU (通过 NIC)
Executor: NIC
示例：
1 1 G0 N0.1 G1 // NIC0 到 NIC1
6.2 多源多目标模式
数据聚合（多源单目标）
Transfer: C0 + G0 -> G1
操作：dst = src_cpu + src_gpu
示例：
1 1 C0 G0 G1 G1 // 两个源，一个目标
数据广播（单源多目标）
Transfer: G0 -> C0 + G1
操作：dst_cpu = dst_gpu = src
示例：
1 1 G0 G0 C0 G1 // 一个源，两个目标
复杂模式（多源多目标）
Transfer: C0 + G0 + G1 -> C1 + G2 + G3
操作：sum = C0 + G0 + G1
C1 = G2 = G3 = sum
示例：
1 1 C0 G0 G1 G0 C1 G2 G3 // 三源三目标
6.3 并发测试
并行传输：
# 4 个并行传输
4 1 C0 G0 G0 C1 G1 G1 G0 G2 G2 G1 G3 G3
执行：所有传输同时执行，测量总带宽
流水线测试：
# H2D + 计算 + D2H 流水线模拟
3 1 C0 D0 G0 N N G0 G0 D0 C0
6.4 细粒度控制
XCC 指定
# 在 GPU0 的 XCC1 上执行
1 1 G0 G0.1 G1
# XCC 亲和性表
gfx.prefXccTable = [[0, 1], [1, 0]]
SDMA 引擎指定
# 使用 GPU0 的 SDMA 引擎 0
1 1 G0 D0.0 C0
# 使用 GPU0 的 SDMA 引擎 1
1 1 G0 D0.1 C0
CU 掩码
// 只使用 CU 0, 2, 4, 6
cfg.gfx.cuMask = {0x55, 0x00, 0x00, 0x00}; // 二进制: 01010101
7. 配置选项详解
7.1 通用选项 (General Options)
struct GeneralOptions {
int numIterations; // 迭代次数（负数表示运行秒数）
int numSubIterations; // 每次迭代的子迭代次数
int numWarmups; // 预热迭代次数
int recordPerIteration; // 记录每次迭代的详细信息
int useInteractive; // 交互模式（暂停等待用户输入）
};
// 默认值
numIterations = 10;
numSubIterations = 1;
numWarmups = 3;
使用示例：
// 运行 100 次迭代
cfg.general.numIterations = 100;
// 运行 60 秒
cfg.general.numIterations = -60;
// 每次迭代重复 10 次
cfg.general.numSubIterations = 10;
7.2 数据选项 (Data Options)
struct DataOptions {
int alwaysValidate; // 每次迭代后验证
int blockBytes; // 数据块大小（字节）
int byteOffset; // 内存分配偏移
vector<float> fillPattern; // 源数据填充模式
int validateDirect; // 直接在 GPU 上验证
int validateSource; // 验证源数据
};
// 默认值
blockBytes = 256;
byteOffset = 0;
7.3 GFX 选项 (GFX Options)
struct GfxOptions {
int blockSize; // 线程块大小
vector<uint32_t> cuMask; // CU 掩码
vector<vector<int>> prefXccTable; // XCC 亲和性表
int unrollFactor; // 展开因子
int useHipEvents; // 使用 HIP Events 计时
int useMultiStream; // 使用多流
int useSingleTeam; // 单团队模式
int waveOrder; // Wavefront 排序
};
// 默认值
blockSize = 256;
unrollFactor = 4;
useHipEvents = 1;
useMultiStream = 0;
useSingleTeam = 0;
waveOrder = 0;
// 约束
blockSize ∈ {64, 128, 192, 256, 320, 384, 448, 512}
unrollFactor ∈ {1, 2, 3, 4, 5, 6, 7, 8}
waveOrder ∈ {0, 1, 2, 3, 4, 5}
7.4 DMA 选项 (DMA Options)
struct DmaOptions {
int useHipEvents; // 使用 HIP Events 计时
int useHsaCopy; // 使用 HSA copy（而非 HIP copy）
};
// 默认值
useHipEvents = 1;
useHsaCopy = 0;
7.5 NIC 选项 (NIC Options)
struct NicOptions {
vector<int> closestNics; // 每个 GPU 的最近 NIC 列表
int ibGidIndex; // GID 索引（-1 为自动）
uint8_t ibPort; // IB 端口号
int ipAddressFamily; // IP 地址族（4 或 6）
int maxRecvWorkReq; // 最大接收工作请求数
int maxSendWorkReq; // 最大发送工作请求数
int queueSize; // 完成队列大小
int roceVersion; // RoCE 版本（1 或 2）
int useRelaxedOrder; // 使用松弛排序
int useNuma; // 切换到最近的 NUMA 节点
};
// 默认值
ibPort = 1;
ipAddressFamily = 4;
maxRecvWorkReq = 16;
maxSendWorkReq = 16;
queueSize = 100;
roceVersion = 2;
useRelaxedOrder = 1;
useNuma = 0;
八、结果输出格式
8.1 测试结果结构
struct TestResults {
int numTimedIterations; // 计时迭代次数
size_t totalBytesTransferred; // 总传输字节数
double avgTotalDurationMsec; // 平均总持续时间
double avgTotalBandwidthGbPerSec; // 平均总带宽
double overheadMsec; // 开销时间
map<ExeDevice, ExeResult> exeResults; // 每执行器结果
vector<TransferResult> tfrResults; // 每传输结果
vector<ErrResult> errResults; // 错误列表
};
8.2 执行器结果
struct ExeResult {
size_t numBytes; // 总字节数
double avgDurationMsec; // 平均持续时间
double avgBandwidthGbPerSec; // 平均带宽
double sumBandwidthGbPerSec; // 传输带宽总和
vector<int> transferIdx; // 传输索引列表
};
输出示例：
Executor: GPU 0 (GFX)
Bandwidth: 125.50 GB/s
Duration: 79.68 ms
Total Bytes: 10000000000
Sum Bandwidth: 130.20 GB/s
Transfers: [0, 1, 2]
8.3 传输结果
struct TransferResult {
size_t numBytes; // 字节数
double avgDurationMsec; // 平均持续时间
double avgBandwidthGbPerSec; // 平均带宽
vector<double> perIterMsec; // 每次迭代时间
vector<set<pair<int,int>>> perIterCUs; // 每次迭代使用的 CU
ExeDevice exeDevice; // 执行器
ExeDevice exeDstDevice; // 目标执行器（NIC）
};
输出示例：
Transfer 0: C0 -> G0 (GPU_GFX 0)
Bandwidth: 42.50 GB/s
Duration: 23.53 ms
Bytes: 1000000000
Per Iteration:
Iter 000: 42.30 GB/s, 23.64 ms, CUs: (0:0,0:1,0:2,0:3)
Iter 001: 42.70 GB/s, 23.42 ms, CUs: (0:0,0:1,0:2,0:3)
...
8.4 错误结果
enum ErrType {
ERR_NONE = 0, // 无错误
ERR_WARN = 1, // 警告
ERR_FATAL = 2, // 致命错误
};
struct ErrResult {
ErrType errType; // 错误类型
std::string errMsg; // 错误消息
};
示例：
WARN: Large BAR not enabled for GPU 1. Multi-GPU access may be limited.
WARN: GPU 0 requests 128 total CUs however only 110 available. Serialization will occur.
FATAL: Transfer 3: Unexpected mismatch at index 1024: Expected 31.00000 Actual: 0.00000
九、高级特性
9.1 硬件拓扑检测
PCIe 树构建：
// 自动检测系统 PCIe 拓扑
GetPCIeTreeRoot();
// 示例拓扑输出
pci0000:00
├── 0000:00:01.0 (GPU 0)
├── 0000:00:02.0 (GPU 1)
├── 0000:00:03.0 (NIC 0)
└── 0000:00:04.0
└── 0000:01:00.0 (GPU 2)
最近设备检测：
// 查找离 GPU 最近的 NUMA 节点
numaNode = GetClosestCpuNumaToGpu(gpuIndex);
// 查找离 GPU 最近的 NIC
nicIndex = GetClosestNicToGpu(gpuIndex);
// 基于 PCIe 树的最低公共祖先 (LCA) 算法
// 考虑总线号距离
9.2 环境变量支持
# GPU 最大硬件队列数
export GPU_MAX_HW_QUEUES=8
# 禁用 SDMA
export HSA_ENABLE_SDMA=0
# NUMA 策略
# （通过 libnuma 自动处理）
9.3 参数验证
配置验证：
ConfigOptionsHaveErrors(cfg, errors);
检查：
- numWarmups >= 0
- blockBytes 是 4 的倍数
- blockSize 是 64 的倍数且 <= MAX_BLOCKSIZE
- unrollFactor <= MAX_UNROLL
- waveOrder < 6
- XCC 表维度正确
- NIC 列表与 GPU 数量匹配
- Large BAR 启用状态
传输验证：
TransfersHaveErrors(cfg, transfers, errors);
检查：
- numBytes > 0
- 至少一个源或目标
- 执行器索引有效
- 内存设备索引有效
- 子执行器数量合理
- SDMA 引擎可用性
- P2P 访问可用性
9.4 交互模式
if (useInteractive) {
printf("Memory prepared:\n");
// 显示所有内存地址
printf("Hit <Enter> to continue: ");
scanf("%*c");
// 执行传输
printf("Transfers complete. Hit <Enter> to continue: ");
scanf("%*c");
}
用途：
- 调试内存分配
- 使用外部工具（如 rocprof）
- 手动验证硬件状态
十、实际应用示例
10.1 基本带宽测试
H2D 带宽：
# 命令行格式
rocm_bandwidth_test "1 16 C0 G0 G0"
解释：
- 1 个传输
- 16 个子执行器（线程块）
- C0: CPU NUMA 节点 0
- G0 (第一个): GPU 0 执行
- G0 (第二个): 写入 GPU 0
D2H 带宽：
rocm_bandwidth_test "1 16 G0 G0 C0"
D2D 带宽：
rocm_bandwidth_test "1 16 G0 G0 G1"
10.2 多 GPU 测试
所有 GPU 对之间的带宽：
# 4 GPU 系统：16 个传输（4x4）
rocm_bandwidth_test "16 8 \
G0 G0 G0 G0 G0 G1 G0 G0 G2 G0 G0 G3 \
G1 G1 G0 G1 G1 G1 G1 G1 G2 G1 G1 G3 \
G2 G2 G0 G2 G2 G1 G2 G2 G2 G2 G2 G3 \
G3 G3 G0 G3 G3 G1 G3 G3 G2 G3 G3 G3"
10.3 高级模式测试
指定传输大小和子执行器：
# -N 表示高级模式
rocm_bandwidth_test "-4 \
C0 G0 G0 16 1G \
G0 G0 C0 16 1G \
G0 G0 G1 8 2G \
G1 G1 G0 8 2G"
解释：
- 4 个传输（高级模式）
- 每个传输指定：src exe dst numCUs size
10.4 配置文件示例
JSON 配置：
{
"general": {
"numIterations": 100,
"numWarmups": 5,
"recordPerIteration": 1
},
"data": {
"blockBytes": 256,
"validateDirect": 1
},
"gfx": {
"blockSize": 256,
"unrollFactor": 4,
"useMultiStream": 1,
"waveOrder": 0
}
}
十一、性能调优指南
11.1 参数选择建议
blockSize 选择：
- 小数据量：128 或 192
- 中等数据量：256（推荐）
- 大数据量：384 或 512
- 考虑寄存器压力
unrollFactor 选择：
- 顺序访问：4-8
- 随机访问：1-2
- 平衡选择：4（默认）
waveOrder 选择：
- 连续内存：0 (U,W,C)
- 分散内存：1 (U,C,W)
- 负载均衡：3 (W,C,U)
numSubExecs 选择：
- CPU: = NUMA 节点上的核心数
- GPU: = CU 数量（或更少以避免序列化）
- 避免过度订阅
11.2 常见性能问题
问题 1：带宽低于预期
可能原因：
1. Large BAR 未启用
2. P2P 不可用
3. 内存未在正确的 NUMA 节点
4. SDMA 被禁用
解决方案：
- 检查 BIOS 设置
- 验证 lspci 输出
- 使用 numactl 验证
- 检查 HSA_ENABLE_SDMA
问题 2：GPU 利用率低
可能原因：
1. 子执行器太少
2. 数据量太小
3. 未使用多流
解决方案：
- 增加 numSubExecs
- 增加传输大小
- 启用 useMultiStream
问题 3：验证失败
可能原因：
1. 内存损坏
2. P2P 配置错误
3. 竞态条件
解决方案：
- 运行内存测试
- 检查 P2P 拓扑
- 增加同步点
11.3 性能分析工具集成
ROCProfiler：
# 启用交互模式
rocm_bandwidth_test --interactive "..."
# 在另一个终端
rocprof --hip-trace ./rocm_bandwidth_test ...
ROCTracer：
export HSA_TOOLS_LIB=/opt/rocm/lib/libroctracer64.so
rocm_bandwidth_test "..."
十二、总结
核心测试原理
ROCm Bandwidth Test (RBT-NG) 通过以下核心机制实现全面的带宽性能测试：
- Transfer 抽象：统一的数据传输模型，支持多源多目标
- 多执行器架构：CPU、GPU GFX、GPU DMA、NIC 四种执行方式
- 精确计时：GPU 硬件时间戳和 HIP Events 双模式
- 内存灵活性：7 种内存类型覆盖所有使用场景
- 并行优化：向量化、循环展开、流并行等技术
- 拓扑感知：PCIe 树分析，最优设备选择
技术特点
- 高精度：纳秒级 GPU 时间戳，准确测量带宽
- 高灵活性：丰富的配置选项适应各种场景
- 高可靠性：完善的验证机制确保数据正确性
- 高性能：多种优化技术最大化硬件利用率
- 可扩展性：插件架构易于添加新功能
适用场景
- 系统性能评估：CPU-GPU 互连带宽基准测试
- 硬件验证：P2P、Large BAR、NUMA 配置验证
- 应用优化：找出数据传输瓶颈
- 多 GPU 调优：评估 GPU 间通信性能
- 网络性能：RDMA 和 RoCE 带宽测试
最佳实践
- 始终启用 Large BAR
- 验证 P2P 拓扑
- 使用多次迭代获得稳定结果
- 根据数据特征选择参数
- 监控 GPU 利用率和温度
- 使用 NUMA 绑定优化 CPU 性能
- 结合 profiling 工具深入分析
扫一扫
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
