rocm_bandwidth_test-6: TB传输完成检测机制详细分析

原创已于 2025-11-29 13:58:19 修改 · 718 阅读

19 ·

CC 4.0 BY-SA版权

文章标签：

#rocm #bandwidth

于 2025-11-28 06:45:00 首次发布

深探ROCm 专栏收录该内容

8 篇文章

订阅专栏

📋 概述

TransferBench 针对不同类型的执行器（CPU、GPU GFX、GPU DMA、NIC）采用了不同的传输完成检测机制。本文档详细分析每种执行器的完成检测策略、计时方法和同步机制。

🎯 核心检测机制分类

1. CPU 执行器 - 基于 std::thread::join()

2. GPU GFX 执行器 - 基于 HIP Events + Stream Sync + 内核时间戳

3. GPU DMA 执行器 - 基于 HIP Events / HSA Signals

4. NIC 执行器 - 基于 InfiniBand 完成队列轮询

1️⃣ CPU 执行器完成检测

检测策略

使用 std::thread::join() 等待所有子线程完成，配合 std::chrono 进行 CPU 时间计时。

代码实现

// 执行单个 CPU Transfer
static void ExecuteCpuTransfer(int const iteration, 
                               ConfigOptions const& cfg, 
                               int const exeIndex, 
                               TransferResources& rss)
{
    // 1️⃣ 启动 CPU 计时器
    auto cpuStart = std::chrono::high_resolution_clock::now();
    
    vector<std::thread> childThreads;

    // 2️⃣ 为每个子执行器（CPU 线程）创建线程
    for (auto const& subExecParam : rss.subExecParamCpu)
        childThreads.emplace_back(std::thread(
            CpuReduceKernel, 
            std::cref(subExecParam), 
            cfg.general.numSubIterations
        ));

    // 3️⃣ 等待所有子线程完成（关键同步点）
    for (auto& subExecThread : childThreads)
        subExecThread.join();  // 🔑 阻塞直到线程完成
    
    childThreads.clear();

    // 4️⃣ 计算耗时
    auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart;
    double deltaMsec = (std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count() * 1000.0) 
                       / cfg.general.numSubIterations;

    // 5️⃣ 记录结果（仅在计时迭代中）
    if (iteration >= 0) {
        rss.totalDurationMsec += deltaMsec;
        if (cfg.general.recordPerIteration)
            rss.perIterMsec.push_back(deltaMsec);
    }
}

执行流程

主线程启动
    ↓
cpuStart = now()
    ↓
创建 N 个子线程
    ├─ Thread 1: CpuReduceKernel(subExecParam[0])
    ├─ Thread 2: CpuReduceKernel(subExecParam[1])
    └─ Thread N: CpuReduceKernel(subExecParam[N-1])
    ↓
for each thread:
    thread.join()  ← 🔑 阻塞等待线程完成
    ↓
所有线程完成
    ↓
cpuDelta = now() - cpuStart
    ↓
记录时间

CPU 核心工作函数

static void CpuReduceKernel(SubExecParam const& p, int numSubIterations)
{
    // 执行 numSubIterations 次子迭代
    for (int iter = 0; iter < numSubIterations; iter++) {
        // 读取源数据并累加
        for (size_t i = 0; i < p.N; i++) {
            float val = 0.0f;
            for (int s = 0; s < p.numSrcs; s++)
                val += p.src[s][i];
            
            // 写入所有目标
            for (int d = 0; d < p.numDsts; d++)
                p.dst[d][i] = val;
        }
    }
    // 函数返回 = 线程完成
}

关键点

✅ 同步机制： std::thread::join() - 阻塞直到线程执行完成
✅ 计时方式： CPU 墙钟时间（wall-clock time）
✅ 优点： 简单可靠，无需硬件支持
✅ 缺点： 包含线程创建和销毁开销

2️⃣ GPU GFX 执行器完成检测

检测策略

采用三重检测机制：

HIP Events - GPU 硬件时间戳（可选，更精确）
hipStreamSynchronize() - 流同步确保完成
内核内部时间戳 - wall_clock64 寄存器读取（用于单个 Transfer 计时）

代码实现

方式 1：单流模式（合并所有 Transfer）

static ErrResult RunGpuExecutor(int const iteration, 
                                ConfigOptions const& cfg, 
                                int const exeIndex, 
                                ExeInfo& exeInfo)
{
    // 1️⃣ CPU 计时开始
    auto cpuStart = std::chrono::high_resolution_clock::now();
    ERR_CHECK(hipSetDevice(exeIndex));

    int numSubExecs = exeInfo.totalSubExecs;
    dim3 const gridSize(xccDim, numSubExecs, 1);
    dim3 const blockSize(cfg.gfx.blockSize, 1);
    hipStream_t stream = exeInfo.streams[0];

    // 2️⃣ 记录 GPU 开始事件（可选）
    if (cfg.gfx.useHipEvents)
        ERR_CHECK(hipEventRecord(exeInfo.startEvents[0], stream));

    // 3️⃣ 启动内核
#if defined(__NVCC__)
    gpuKernel<<<gridSize, blockSize, 0, stream>>>(
        exeInfo.subExecParamGpu, 
        cfg.gfx.waveOrder, 
        cfg.general.numSubIterations
    );
#else
    hipExtLaunchKernelGGL(
        gpuKernel,
        gridSize, blockSize,
        0,  // sharedMem
        stream,
        cfg.gfx.useHipEvents ? exeInfo.startEvents[0] : NULL,  // 开始事件
        cfg.gfx.useHipEvents ? exeInfo.stopEvents[0] : NULL,   // 结束事件
        0,  // flags
        exeInfo.subExecParamGpu,
        cfg.gfx.waveOrder,
        cfg.general.numSubIterations
    );
#endif

    // 4️⃣ 记录 GPU 结束事件（可选）
    if (cfg.gfx.useHipEvents)
        ERR_CHECK(hipEventRecord(exeInfo.stopEvents[0], stream));

    // 5️⃣ 🔑 关键：同步流，确保内核完成
    ERR_CHECK(hipStreamSynchronize(stream));

    // 6️⃣ CPU 计时结束
    auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart;
    double cpuDeltaMsec = std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count() 
                          * 1000.0 / cfg.general.numSubIterations;

    // 7️⃣ 选择计时方式
    if (iteration >= 0) {
        if (cfg.gfx.useHipEvents) {
            // 使用 GPU 硬件时间（更精确）
            float gpuDeltaMsec;
            ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, 
                                          exeInfo.startEvents[0], 
                                          exeInfo.stopEvents[0]));
            gpuDeltaMsec /= cfg.general.numSubIterations;
            exeInfo.totalDurationMsec += gpuDeltaMsec;
        } else {
            // 使用 CPU 计时
            exeInfo.totalDurationMsec += cpuDeltaMsec;
        }

        // 8️⃣ 从内核读取每个 Transfer 的时间（使用 wall_clock64）
        for (int i = 0; i < exeInfo.resources.size(); i++) {
            TransferResources& rss = exeInfo.resources[i];
            long long minStartCycle = std::numeric_limits<long long>::max();
            long long maxStopCycle = std::numeric_limits<long long>::min();

            // 找到这个 Transfer 所有子执行器的最早开始和最晚结束时间
            for (auto subExecIdx : rss.subExecIdx) {
                minStartCycle = std::min(minStartCycle, 
                    exeInfo.subExecParamGpu[subExecIdx].startCycle);
                maxStopCycle = std::max(maxStopCycle, 
                    exeInfo.subExecParamGpu[subExecIdx].stopCycle);
            }
            
            // 计算 Transfer 耗时
            double deltaMsec = (maxStopCycle - minStartCycle) 
                             / (double)(exeInfo.wallClockRate);
            deltaMsec /= cfg.general.numSubIterations;
            rss.totalDurationMsec += deltaMsec;
        }
    }
    return ERR_NONE;
}

方式 2：多流模式（每个 Transfer 独立流）

if (cfg.gfx.useMultiStream) {
    // 每个 Transfer 使用独立的流和事件
    vector<std::future<ErrResult>> asyncTransfers;
    
    for (int i = 0; i < exeInfo.streams.size(); i++) {
        asyncTransfers.emplace_back(std::async(
            std::launch::async,
            ExecuteGpuTransfer,
            iteration,
            exeInfo.streams[i],
            cfg.gfx.useHipEvents ? exeInfo.startEvents[i] : NULL,
            cfg.gfx.useHipEvents ? exeInfo.stopEvents[i] : NULL,
            xccDim,
            std::cref(cfg),
            std::ref(exeInfo.resources[i])
        ));
    }
    
    // 🔑 等待所有异步 Transfer 完成
    for (auto& asyncTransfer : asyncTransfers)
        ERR_CHECK(asyncTransfer.get());
}

GPU 内核内部时间戳

template<typename PACKED_FLOAT, int BLOCKSIZE, int UNROLL, int TEMPORAL_MODE>
__global__ void __launch_bounds__(BLOCKSIZE) 
GpuReduceKernel(SubExecParam* params, int waveOrder, int numSubIterations)
{
    // 1️⃣ 线程 0 记录开始时间戳
    int64_t startCycle;
    if (threadIdx.x == 0)
        startCycle = GetTimestamp();  // wall_clock64 或 clock64

    SubExecParam& p = params[blockIdx.y];

    // ... 执行数据传输和计算 ...

    int subIterations = 0;
    while (1) {
        // 三层循环处理数据
        // Loop 1: 主要的展开循环
        // Loop 2: 处理剩余的 PACKED_FLOAT
        // Loop 3: 处理剩余的单精度浮点数
        
        if (++subIterations == numSubIterations)
            break;
    }

    // 2️⃣ 同步所有线程
    __syncthreads();
    
    // 3️⃣ 线程 0 记录结束时间戳和硬件信息
    if (threadIdx.x == 0) {
        __threadfence_system();  // 🔑 确保所有内存写入全局可见
        p.stopCycle = GetTimestamp();
        p.startCycle = startCycle;
        GetHwId(p.hwId);    // 获取 CU/SM ID
        GetXccId(p.xccId);  // 获取 XCC ID
    }
}

时间戳函数

// 设备端时间戳获取
__device__ int64_t GetTimestamp()
{
#if defined(__NVCC__)
    int64_t time;
    asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(time));
    return time;
#else
    return wall_clock64();  // AMD ROCm 内置函数
#endif
}

同步机制层次

Level 1: hipStreamSynchronize(stream)
    ↓ 确保 GPU 内核执行完成
    
Level 2: __syncthreads() (内核内)
    ↓ 确保线程块内所有线程完成
    
Level 3: __threadfence_system() (内核内)
    ↓ 确保所有内存写入对系统可见
    
Level 4: hipEventElapsedTime()
    ↓ 计算 GPU 硬件精确时间

关键点

✅ 硬件事件： hipEventRecord() 在 GPU 硬件级别记录时间戳
✅ 流同步： hipStreamSynchronize() 阻塞直到流中所有操作完成
✅ 内核时间戳： wall_clock64() 提供纳秒级精度
✅ 内存栅栏： __threadfence_system() 确保数据可见性
✅ 优点： 精确度高，硬件级计时，开销小
✅ 缺点： 需要 GPU 硬件支持

3️⃣ GPU DMA 执行器完成检测

检测策略

支持两种方式：

HIP DMA - 使用 hipMemcpyAsync() + hipStreamSynchronize() + HIP Events
HSA DMA - 使用 hsa_amd_memory_async_copy() + HSA Signals

代码实现

方式 1：HIP DMA

static ErrResult ExecuteDmaTransfer(int const iteration,
                                    bool const useSubIndices,
                                    hipStream_t const stream,
                                    hipEvent_t const startEvent,
                                    hipEvent_t const stopEvent,
                                    ConfigOptions const& cfg,
                                    TransferResources& resources)
{
    // 1️⃣ CPU 计时开始
    auto cpuStart = std::chrono::high_resolution_clock::now();

    int subIterations = 0;
    if (!useSubIndices && !cfg.dma.useHsaCopy) {
        // === HIP DMA 路径 ===
        
        // 2️⃣ 记录开始事件
        if (cfg.dma.useHipEvents)
            ERR_CHECK(hipEventRecord(startEvent, stream));

        // 3️⃣ 执行异步拷贝
        do {
            ERR_CHECK(hipMemcpyAsync(
                resources.dstMem[0], 
                resources.srcMem[0], 
                resources.numBytes, 
                hipMemcpyDefault, 
                stream
            ));
        } while (++subIterations != cfg.general.numSubIterations);

        // 4️⃣ 记录结束事件
        if (cfg.dma.useHipEvents)
            ERR_CHECK(hipEventRecord(stopEvent, stream));
        
        // 5️⃣ 🔑 同步流，确保 DMA 完成
        ERR_CHECK(hipStreamSynchronize(stream));
    }
    
    // ... HSA 路径见下方 ...

    // 6️⃣ CPU 计时结束
    auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart;
    double cpuDeltaMsec = std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count() 
                          * 1000.0 / cfg.general.numSubIterations;

    // 7️⃣ 记录结果
    if (iteration >= 0) {
        double deltaMsec = cpuDeltaMsec;
        if (!useSubIndices && !cfg.dma.useHsaCopy && cfg.dma.useHipEvents) {
            // 使用 GPU 事件时间
            float gpuDeltaMsec;
            ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, startEvent, stopEvent));
            deltaMsec = gpuDeltaMsec / cfg.general.numSubIterations;
        }
        resources.totalDurationMsec += deltaMsec;
        if (cfg.general.recordPerIteration)
            resources.perIterMsec.push_back(deltaMsec);
    }
    return ERR_NONE;
}

方式 2：HSA DMA（更底层，支持指定 SDMA 引擎）

else {
    // === HSA DMA 路径 ===
    do {
        // 1️⃣ 初始化信号（值设为 1）
        hsa_signal_store_screlease(resources.signal, 1);
        
        if (!useSubIndices) {
            // 2️⃣ 标准 HSA 异步拷贝
            ERR_CHECK(hsa_amd_memory_async_copy(
                resources.dstMem[0],
                resources.dstAgent,
                resources.srcMem[0],
                resources.srcAgent,
                resources.numBytes,
                0,         // num_dep_signals
                NULL,      // dep_signals
                resources.signal  // 完成信号
            ));
        } else {
            // 3️⃣ 指定 SDMA 引擎的拷贝
            HSA_CALL(hsa_amd_memory_async_copy_on_engine(
                resources.dstMem[0],
                resources.dstAgent,
                resources.srcMem[0],
                resources.srcAgent,
                resources.numBytes,
                0,
                NULL,
                resources.signal,
                resources.sdmaEngineId,  // 指定引擎
                true
            ));
        }
        
        // 4️⃣ 🔑 等待信号变为 0（轮询等待 DMA 完成）
        while (hsa_signal_wait_scacquire(
                   resources.signal, 
                   HSA_SIGNAL_CONDITION_LT,  // 小于
                   1,                        // 阈值
                   UINT64_MAX,               // 超时
                   HSA_WAIT_STATE_ACTIVE     // 活动等待
               ) >= 1)
            ;  // 🔄 忙等待直到信号 < 1
            
    } while (++subIterations != cfg.general.numSubIterations);
}

HSA Signal 工作原理

初始状态:
    signal = 1  (store_screlease)
    ↓
启动 DMA:
    hsa_amd_memory_async_copy(..., signal)
    ↓
DMA 引擎执行拷贝
    ↓
DMA 完成时:
    signal 自动递减为 0
    ↓
CPU 轮询检测:
    while (signal_wait(...) >= 1)
        ↓
    signal < 1 → 退出循环 ✓ 完成

DMA 执行器主循环

static ErrResult RunDmaExecutor(int const iteration, 
                                ConfigOptions const& cfg, 
                                int const exeIndex, 
                                ExeInfo& exeInfo)
{
    auto cpuStart = std::chrono::high_resolution_clock::now();
    ERR_CHECK(hipSetDevice(exeIndex));

    // 🔑 并行启动所有 Transfer（每个 Transfer 独立流）
    vector<std::future<ErrResult>> asyncTransfers;
    for (int i = 0; i < exeInfo.resources.size(); i++) {
        asyncTransfers.emplace_back(std::async(
            std::launch::async,
            ExecuteDmaTransfer,
            iteration,
            exeInfo.useSubIndices,
            exeInfo.streams[i],
            cfg.dma.useHipEvents ? exeInfo.startEvents[i] : NULL,
            cfg.dma.useHipEvents ? exeInfo.stopEvents[i] : NULL,
            std::cref(cfg),
            std::ref(exeInfo.resources[i])
        ));
    }

    // 等待所有异步 Transfer 完成
    for (auto& asyncTransfer : asyncTransfers)
        ERR_CHECK(asyncTransfer.get());

    // 记录总时间
    auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart;
    double deltaMsec = std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count() 
                       * 1000.0 / cfg.general.numSubIterations;
    if (iteration >= 0)
        exeInfo.totalDurationMsec += deltaMsec;
    
    return ERR_NONE;
}

关键点

✅ HIP 路径： hipStreamSynchronize() 阻塞等待
✅ HSA 路径： hsa_signal_wait_scacquire() 轮询等待
✅ 信号机制： HSA Signal 由硬件自动更新，无需软件干预
✅ 引擎选择： 可指定特定 SDMA 引擎（多引擎并行）
✅ 优点： HSA 路径开销更小，支持更细粒度控制
✅ 缺点： HSA API 仅 AMD 平台支持

4️⃣ NIC (InfiniBand/RDMA) 执行器完成检测

检测策略

使用 InfiniBand Verbs 完成队列轮询 - ibv_poll_cq()

代码实现

static ErrResult RunNicExecutor(int const iteration, 
                                ConfigOptions const& cfg, 
                                int const exeIndex, 
                                ExeInfo& exeInfo)
{
    // NUMA 绑定
    if (cfg.nic.useNuma) {
        int numaNode = GetIbvDeviceList()[exeIndex].numaNode;
        if (numaNode != -1)
            numa_run_on_node(numaNode);
    }

    auto transferCount = exeInfo.resources.size();
    std::vector<double> totalTimeMsec(transferCount, 0.0);

    int subIterations = 0;
    auto cpuStart = std::chrono::high_resolution_clock::now();
    std::vector<std::chrono::high_resolution_clock::time_point> transferTimers(transferCount);

    do {
        // 1️⃣ 跟踪每个 Transfer 的 QP 完成数量
        std::vector<uint8_t> receivedQPs(transferCount, 0);
        
        // 2️⃣ 发起所有 Transfer 的 RDMA 发送
        for (auto i = 0; i < transferCount; i++) {
            transferTimers[i] = std::chrono::high_resolution_clock::now();
            ERR_CHECK(ExecuteNicTransfer(iteration, cfg, exeIndex, exeInfo.resources[i]));
        }
        
        // 3️⃣ 🔑 轮询完成队列，直到所有 Transfer 完成
        size_t completedTransfers = 0;
        while (completedTransfers < transferCount) {
            for (auto i = 0; i < transferCount; i++) {
                if (receivedQPs[i] < exeInfo.resources[i].qpCount) {
                    auto& rss = exeInfo.resources[i];
                    
                    // 4️⃣ 轮询完成队列
                    ibv_wc wc;  // Work Completion
                    int nc = ibv_poll_cq(rss.srcCompQueue, 1, &wc);
                    
                    if (nc > 0) {
                        // ✅ 收到完成通知
                        receivedQPs[i]++;
                        
                        // 检查状态
                        if (wc.status != IBV_WC_SUCCESS) {
                            return {ERR_FATAL, "Transfer %d: Received unsuccessful work completion", 
                                    rss.transferIdx};
                        }
                    } else if (nc < 0) {
                        // ❌ 轮询错误
                        return {ERR_FATAL, "Transfer %d: Received negative work completion", 
                                rss.transferIdx};
                    }
                    // nc == 0: 暂无完成，继续轮询
                    
                    // 5️⃣ 所有 QP 完成，记录时间
                    if (receivedQPs[i] == rss.qpCount) {
                        auto cpuDelta = std::chrono::high_resolution_clock::now() - transferTimers[i];
                        double deltaMsec = std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count() * 1000.0;
                        
                        if (iteration >= 0) {
                            totalTimeMsec[i] += deltaMsec;
                        }
                        completedTransfers++;
                    }
                }
            }
        }  // 🔄 继续轮询未完成的 Transfer
        
    } while (++subIterations < cfg.general.numSubIterations);

    // 6️⃣ 记录总时间
    auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart;
    double deltaMsec = std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count() 
                       * 1000.0 / cfg.general.numSubIterations;

    if (iteration >= 0) {
        exeInfo.totalDurationMsec += deltaMsec;
        for (int i = 0; i < transferCount; i++) {
            auto& rss = exeInfo.resources[i];
            double transferTimeMsec = totalTimeMsec[i] / cfg.general.numSubIterations;
            rss.totalDurationMsec += transferTimeMsec;
            if (cfg.general.recordPerIteration)
                rss.perIterMsec.push_back(transferTimeMsec);
        }
    }
    return ERR_NONE;
}

NIC Transfer 发送

static ErrResult ExecuteNicTransfer(int const iteration, 
                                    ConfigOptions const& cfg, 
                                    int const exeIndex, 
                                    TransferResources& rss)
{
    // 为每个 Queue Pair 发送工作请求
    ibv_send_wr* badWorkReq;
    for (int qpIndex = 0; qpIndex < rss.qpCount; qpIndex++) {
        int error = ibv_post_send(
            rss.srcQueuePairs[qpIndex], 
            &rss.sendWorkRequests[qpIndex], 
            &badWorkReq
        );
        if (error)
            return {ERR_FATAL, "Transfer %d: Error when calling ibv_post_send for QP %d Error code %d\n",
                    rss.transferIdx, qpIndex, error};
    }
    return ERR_NONE;
}

InfiniBand 完成检测流程

发起 RDMA 操作:
    ibv_post_send(QP, work_request)
    ↓
    work_request 提交到 NIC 硬件
    ↓
NIC 执行 RDMA 传输
    ↓
传输完成:
    NIC 在完成队列 (CQ) 中生成 Work Completion (WC)
    ↓
CPU 轮询检测:
    loop:
        nc = ibv_poll_cq(CQ, 1, &wc)
        if nc > 0:
            ✓ 收到完成
            检查 wc.status
            break
        else if nc == 0:
            ← 继续轮询
        else:
            ✗ 错误

关键点

✅ 轮询机制： ibv_poll_cq() 非阻塞轮询完成队列
✅ 多 QP 支持： 每个 Transfer 可使用多个 Queue Pair 并行
✅ 状态检查： Work Completion 包含状态码验证传输成功
✅ CPU 时间： 使用 CPU 计时，包含轮询开销
✅ 优点： 支持 RDMA 零拷贝，低延迟
✅ 缺点： 轮询消耗 CPU，需要 InfiniBand 硬件

🔄 完成检测机制对比

执行器类型	同步机制	计时方式	精度	开销	硬件要求
CPU	`std::thread::join()`	CPU 墙钟时间	微秒级	低	无
GPU GFX	`hipStreamSynchronize()` + HIP Events + 内核时间戳	GPU 硬件时钟	纳秒级	极低	GPU
GPU DMA (HIP)	`hipStreamSynchronize()` + HIP Events	GPU 硬件时钟	纳秒级	极低	GPU
GPU DMA (HSA)	`hsa_signal_wait()` 轮询	CPU 墙钟时间	微秒级	低	AMD GPU
NIC	`ibv_poll_cq()` 轮询	CPU 墙钟时间	微秒级	中（轮询）	InfiniBand

🎯 最佳实践建议

GPU GFX 执行器

// ✅ 推荐：启用 HIP Events 获得最精确计时
cfg.gfx.useHipEvents = 1;

// ✅ 单流模式：适合多个小 Transfer
cfg.gfx.useMultiStream = 0;

// ✅ 多流模式：适合少量大 Transfer，提高并发
cfg.gfx.useMultiStream = 1;

GPU DMA 执行器

// ✅ HIP 路径：更通用，支持 NVIDIA/AMD
cfg.dma.useHsaCopy = 0;
cfg.dma.useHipEvents = 1;

// ✅ HSA 路径：AMD 专用，更底层，支持引擎选择
cfg.dma.useHsaCopy = 1;  // AMD only

CPU 执行器

// ✅ 绑定 NUMA 节点提高性能
numa_run_on_node(nodeId);

// ✅ 子执行器数量 = CPU 核心数
int numSubExecs = GetNumSubExecutors({EXE_CPU, nodeId});

NIC 执行器

// ✅ 启用 NUMA 绑定
cfg.nic.useNuma = 1;

// ✅ 调整队列大小
cfg.nic.queueSize = 100;
cfg.nic.maxSendWorkReq = 16;

📊 性能开销分析

同步开销排序（从低到高）

GPU GFX (HIP Events) - ~1-2 微秒
- 硬件时间戳，无 CPU 干预
GPU DMA (HSA Signal) - ~5-10 微秒
- 高效轮询，信号由硬件更新
GPU DMA (HIP Events) - ~1-2 微秒
- 与 GFX 类似
CPU (thread::join) - ~10-50 微秒
- 取决于线程调度
NIC (ibv_poll_cq) - ~100-500 微秒
- 持续轮询，消耗 CPU 周期

减少开销的技巧

// 1. GPU: 批量内核启动减少同步次数
cfg.general.numSubIterations = 10;  // 一次启动执行多次

// 2. DMA: 使用 HSA 路径（AMD）
cfg.dma.useHsaCopy = 1;

// 3. NIC: 使用更大的消息减少轮询次数
numBytesPerTransfer = 1 << 20;  // 1 MB

// 4. 所有：减少预热迭代次数
cfg.general.numWarmups = 1;

🔍 调试技巧

启用详细计时

// 记录每次迭代的时间
cfg.general.recordPerIteration = 1;

// 查看 CU 使用情况（仅 GPU GFX）
for (auto& cuPair : results.tfrResults[0].perIterCUs[0]) {
    printf("XCC %d, CU %d\n", cuPair.first, cuPair.second);
}

检测传输失败

// 检查错误结果
if (!RunTransfers(cfg, transfers, results)) {
    for (auto& err : results.errResults) {
        printf("[%s] %s\n", 
               err.errType == ERR_FATAL ? "FATAL" : "WARN", 
               err.errMsg.c_str());
    }
}

验证完成性

// GPU: 检查内核启动后的错误
hipError_t err = hipGetLastError();
if (err != hipSuccess) {
    printf("Kernel launch failed: %s\n", hipGetErrorString(err));
}

// DMA: 检查信号值
printf("Signal value: %ld\n", hsa_signal_load_relaxed(signal));

// NIC: 检查 Work Completion 状态
if (wc.status != IBV_WC_SUCCESS) {
    printf("RDMA error: %s\n", ibv_wc_status_str(wc.status));
}