如何将CUDA事件时间转换到CPU时间轴？

一、将CUDA事件转换为CPU时间轴的绝对时间 #

方法1：基于基准同步点的时间映射 #

通过建立GPU事件时间与CPU时间的基准对应关系，将后续GPU事件的相对时间转换为CPU绝对时间：

// 创建基准事件并同步GPU与CPU时间
cudaEvent_t baseEvent;
cudaEventCreate(&baseEvent);
cudaEventRecord(baseEvent);
cudaDeviceSynchronize();
// 记录CPU基准时间（程序启动以来的绝对时间）
auto cpuBaseTime = std::chrono::high_resolution_clock::now();

// 原有的kernel循环
for(int i{0}; i < n; ++i){
    auto& startEvent = startEvents[i];
    auto& stopEvent = stopEvents[i];

    cudaEventRecord(startEvent);
    kernel<<<1000, dim3{32, 32, 1}>>>();
    cudaEventRecord(stopEvent);
}

cudaDeviceSynchronize();

// 转换每个事件到CPU绝对时间
float msOffset;
// 计算第一个启动事件相对基准事件的时间差
cudaEventElapsedTime(&msOffset, baseEvent, startEvents[0]);
auto startAbsTime = cpuBaseTime + std::chrono::duration<float, std::milli>(msOffset);

// 同理处理其他start/stop事件，通过cudaEventElapsedTime计算相对基准的偏移，叠加到CPU基准时间上

注意：需确保基准事件是在所有前置GPU任务完成后记录的，否则偏移计算会包含前置任务的耗时。

方法2：异步轮询事件完成并记录CPU时间 #

通过非阻塞的cudaEventQuery轮询GPU事件状态，一旦事件完成立即记录当前CPU时间，避免同步等待的开销：

// 启动kernel并记录事件（原循环逻辑不变）
for(int i{0}; i < n; ++i){
    auto& startEvent = startEvents[i];
    auto& stopEvent = stopEvents[i];

    cudaEventRecord(startEvent);
    kernel<<<1000, dim3{32, 32, 1}>>>();
    cudaEventRecord(stopEvent);
}

// 异步轮询所有事件，记录完成时的CPU时间
std::vector<std::chrono::high_resolution_clock::time_point> startAbsTimes(n), stopAbsTimes(n);
std::vector<bool> startDone(n, false), stopDone(n, false);

while(!std::all_of(startDone.begin(), startDone.end(), [](bool b){return b;}) || 
      !std::all_of(stopDone.begin(), stopDone.end(), [](bool b){return b;})){
    for(int i=0; i<n; ++i){
        if(!startDone[i]){
            if(cudaEventQuery(startEvents[i]) == cudaSuccess){
                startAbsTimes[i] = std::chrono::high_resolution_clock::now();
                startDone[i] = true;
            }
        }
        if(!stopDone[i]){
            if(cudaEventQuery(stopEvents[i]) == cudaSuccess){
                stopAbsTimes[i] = std::chrono::high_resolution_clock::now();
                stopDone[i] = true;
            }
        }
    }
    // 短暂休眠，降低CPU占用
    std::this_thread::sleep_for(std::chrono::microseconds(10));
}

优势：无需等待所有GPU任务完成，可异步捕获每个事件完成的CPU时间；不足：轮询会占用少量CPU资源，需合理设置休眠时间平衡精度与开销。

关于cudaLaunchHostFunc的说明 #

cudaLaunchHostFunc并不会阻塞流中的GPU后续任务，它只是将主机函数加入流的任务队列，仅当流中前置GPU任务完成后，才会在主机线程执行该函数。若担心主机线程负载过高，可将回调函数绑定到单独的主机线程（通过创建专用线程处理回调逻辑），避免影响主线程的其他任务。目前CUDA官方未提供cudaLaunchHostFuncAsync这类完全异步的主机回调接口。

二、NSight获取kernel时间戳的原理 #

NSight依赖CUDA Profiler Interface（CUPTI）实现精确时间戳捕获：

• CUPTI通过GPU硬件的性能计数器与事件钩子，直接获取kernel在GPU上的启动、结束硬件时间戳；
• NSight会在程序运行时建立GPU时钟与CPU时钟的同步映射，将GPU硬件时间转换为CPU时间轴的绝对时间；
• 整个过程是低侵入式的，无需在应用代码中添加额外同步逻辑，通过底层硬件与驱动层的支持实现无干扰的性能采样。

内容的提问来源于stack exchange，提问作者nikitablack