如何将 CUDA 时钟周期转换为毫秒？

Posted 2023-04-15

【中文标题】如何将 CUDA 时钟周期转换为毫秒？

【英文标题】：How to convert CUDA clock cycles to milliseconds?

【发布时间】：2017-08-17 22:03:21

【问题描述】：

我想测量一些代码在我的内核所花费的时间。我已经关注了this question 以及它的 cmets，所以我的内核看起来像这样：

__global__ void kernel(..., long long int *runtime)

    long long int start = 0; 
    long long int stop = 0;

    asm volatile("mov.u64 %0, %%clock64;" : "=l"(start));

    /* Some code here */

    asm volatile("mov.u64 %0, %%clock64;" : "=l"(stop));

    runtime[threadIdx.x] = stop - start;
    ...

答案说要进行如下转换：

计时器计算时钟滴答数。要获得毫秒数，请将其除以设备上的 GHz 数，然后乘以 1000。

我这样做：

for(long i = 0; i < size; i++)

  fprintf(stdout, "%d:%ld=%f(ms)\n", i,runtime[i], (runtime[i]/1.62)*1000.0);

其中 1.62 是我设备的 GPU 最大时钟频率。但是我以毫秒为单位的时间看起来并不正确，因为它表明每个线程都需要几分钟才能完成。这不可能是正确的，因为执行在不到一秒的挂钟时间内完成。转换公式不正确还是我在某处犯了错误？谢谢。

【问题讨论】：

除以赫兹数，而不是 GHz。除以1620000000.0f。时钟周期除以每秒的时钟周期得出秒数。将秒数乘以 1000 得到毫秒数。

@RobertCrovella，现在按预期工作，谢谢！如果您将此作为答案发布，我很乐意将其标记为已接受。

【参考方案1】：

在您的情况下正确的转换不是 GHz：

fprintf(stdout, "%d:%ld=%f(ms)\n", i,runtime[i], (runtime[i]/1.62)*1000.0);
                                                             ^^^^

但是赫兹：

fprintf(stdout, "%d:%ld=%f(ms)\n", i,runtime[i], (runtime[i]/1620000000.0f)*1000.0);
                                                             ^^^^^^^^^^^^^

在维度分析中：

                  clock cycles
clock cycles  /  -------------- = seconds
                   second
                    

第一项是时钟周期测量。第二项是 GPU 的频率（以赫兹为单位，而不是 GHz），第三项是所需的测量值（秒）。您可以通过将秒乘以 1000 来转换为毫秒。

这是一个工作示例，展示了一种独立于设备的方法（因此您不必对时钟频率进行硬编码）：

$ cat t1306.cu
#include <stdio.h>

const long long delay_time = 1000000000;
const int nthr = 1;
const int nTPB = 256;

__global__ void kernel(long long *clocks)

  int idx=threadIdx.x+blockDim.x*blockIdx.x;
  long long start=clock64();
  while (clock64() < start+delay_time);
  if (idx < nthr) clocks[idx] = clock64()-start;


int main()

  int peak_clk = 1;
  int device = 0;
  long long *clock_data;
  long long *host_data;
  host_data = (long long *)malloc(nthr*sizeof(long long));
  cudaError_t err = cudaDeviceGetAttribute(&peak_clk, cudaDevAttrClockRate, device);
  if (err != cudaSuccess) printf("cuda err: %d at line %d\n", (int)err, __LINE__); return 1;
  err = cudaMalloc(&clock_data, nthr*sizeof(long long));
  if (err != cudaSuccess) printf("cuda err: %d at line %d\n", (int)err, __LINE__); return 1;
  kernel<<<(nthr+nTPB-1)/nTPB, nTPB>>>(clock_data);
  err = cudaMemcpy(host_data, clock_data, nthr*sizeof(long long), cudaMemcpyDeviceToHost);
  if (err != cudaSuccess) printf("cuda err: %d at line %d\n", (int)err, __LINE__); return 1;
  printf("delay clock cycles: %ld, measured clock cycles: %ld, peak clock rate: %dkHz, elapsed time: %fms\n", delay_time, host_data[0], peak_clk, host_data[0]/(float)peak_clk);
  return 0;

$ nvcc -arch=sm_35 -o t1306 t1306.cu
$ ./t1306
delay clock cycles: 1000000000, measured clock cycles: 1000000210, peak clock rate: 732000kHz, elapsed time: 1366.120483ms
$

这使用cudaDeviceGetAttribute 来获取时钟频率，它返回一个以 kHz 为单位的结果，这使我们能够轻松地计算这种情况下的毫秒数。

根据我的经验，上述方法通常适用于时钟速率以报告速率运行的数据中心 GPU（可能会受到您在 nvidia-smi 中所做的设置的影响。）其他 GPU（例如 GeForce GPU）可能运行在（不可预测的）提升时钟，这会使该方法不准确。

此外，最近，CUDA 能够抢占 GPU 上的活动。这可以在多种情况下发生，例如调试、CUDA 动态并行和其他情况。如果出于某种原因发生抢占，尝试基于clock64() 测量任何内容通常是不可靠的。

【讨论】：

我不太明白延迟时间和while (clock64() < start+delay_time);的重要性。你可以解释吗？谢谢。

我假设它纯粹是模拟由一些实际工作引起的延迟，这个假设是否正确？

是的，它强制内核持续时间至少与delay_time一样长

【参考方案2】：

clock64 以图形时钟周期返回一个值。图形时钟是动态的，所以我不建议使用常量来尝试转换为秒。如果您想转换为挂钟时间，那么更好的选择是使用globaltimer，这是一个 64 位时钟寄存器，可通过以下方式访问：

asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(start));

单位是纳秒。

默认分辨率为 32ns，每 µs 更新一次。 NVIDIA 性能工具强制每 32 ns（或 31.25 MHz）更新一次。 CUPTI 在捕获并发内核跟踪时使用此时钟作为开始时间。