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

Question

我想衡量一段代码的时间内，我的内核负责。我已经关注了这个问题及其评论，以便我的内核看起来像这样：

__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 最大时钟频率。但是我以毫秒为单位的时间看起来不正确，因为它表明每个线程需要几分钟才能完成。这不可能是正确的，因为执行在不到一秒的挂钟时间内完成。转换公式不正确还是我在某处犯了错误？谢谢。

Answer 1

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

\n

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

\n

单位为纳秒。

\n

默认分辨率为 32ns，每 \xc2\xb5s 更新一次。NVIDIA 性能工具强制每 32 ns（或 31.25 MHz）更新一次。当捕获并发内核跟踪时，CUPTI 使用此时钟作为启动时间。

\n

Answer 2

在您的情况下正确的转换不是 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()通常是不可靠的。