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

我想用一些代码来测量时间within我的内核需要。我已经关注了这个问题 https://stackoverflow.com/questions/11209228/timing-different-sections-in-cuda-kernel连同它的注释，这样我的内核看起来像这样：

__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：

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 http://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__DEVICE.html#group__CUDART__DEVICE_1gb22e8256592b836df9a9cc36c9db7151获取时钟速率，它返回以 kHz 为单位的结果，这使我们能够在这种情况下轻松计算毫秒。

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

此外，最近，CUDA 能够抢占 GPU 上的活动。这可以在多种情况下发生，例如调试、CUDA 动态并行性和其他情况。如果由于某种原因发生抢占，则尝试根据clock64()一般是不可靠的。