使用 CUDA 时如何测量每个块的执行时间？答案

【问题标题】：How to measure the execution time of every block when using CUDA?使用 CUDA 时如何测量每个块的执行时间？
【发布时间】：2011-04-03 00:03:27
【问题描述】：

clock() 不够准确。

【问题讨论】：

我相信他的意思是线程块，CUDA 使用的术语。
我的意思是每个块都被分配了整个计算的不同部分。我想测试每个块用于完成自己的任务的准确时间。
我一直在使用 clock() 但没有注意到它有多不准确。你能告诉我吗？我认为它是准确的，因为它调用了 GPU 内部的时钟性能计数器。
@cnhk 您能否在您的问题中提供更多信息，以便该问题对其他 stackoverflow 用户有用。 clock() 设备函数是周期准确的，因此不清楚您为什么认为它不准确。不准确的原因可能包括不正确使用clock()、未能处理翻转、编译器移动SASS指令等。

标签： cuda gpu parallel-processing

【解决方案1】：

使用 CUDA 事件来测量内核时间或 CUDA 操作（memcpy 等）：

// Prepare
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
// Start record
cudaEventRecord(start, 0);
// Do something on GPU
MyKernel<<<dimGrid, dimBlock>>>(input_data, output_data);
// Stop event
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop); // that's our time!
// Clean up:
cudaEventDestroy(start);
cudaEventDestroy(stop);

请参阅 CUDA 编程指南，第 3.2.7.6 节

【讨论】：

这给出了整个内核的总执行时间，而不是我认为的单个块。
当然。调度程序以随机顺序执行块，块的数量很大（您必须有内存来存储所有计时器），以及主要的事情：您到底想测量什么时间？由于 RAM 延迟，大部分时间块都处于睡眠状态，等待内存。但是由于 GPU 的高负载，当一个块在等待时，另一个块正在运行。因此特定块的执行时间可能非常大。我认为平均时间是最重要的特征。

【解决方案2】：

如何在每个 CUDA 线程中使用 clock() 函数来计算开始和结束时间。并将其存储在一个数组中，这样您就可以根据数组索引确定哪个线程在何时开始/停止，如下所示：

__global__ void kclock(unsigned int *ts) {
    unsigned int start_time = 0, stop_time = 0;

    start_time = clock();

    // Code we need to measure should go here.

    stop_time = clock();

    ts[(blockIdx.x * blockDim.x + threadIdx.x) * 2] = start_time;
    ts[(blockIdx.x * blockDim.x + threadIdx.x) * 2 + 1] = stop_time; 
}

然后使用此数组计算您正在考虑的块的最小开始时间和最大停止时间。例如，您可以计算对应于 CUDA 中 (0, 0) 块的时间数组的索引范围，并使用 min/max 来计算执行时间。

【讨论】：

为了减少内存消耗，我建议在扭曲开始时添加条件逻辑，只有 threadIdx.x == 0（假设 1D blockDim）来读取和写入时间戳。在内核结束时，我建议让每个 warp 的 laneIdx == 0 将停止时间写入相同的内存位置。这将是一个竞争条件，但会相当准确。

【解决方案3】：

我认为 long long int clock64() 是你要找的吗？

请参阅 Cuda 编程指南，C 语言扩展，B. 11。

【讨论】：