CUDA编程笔记004. 测量kernel耗时

1、使用cpu时钟测量kernel耗时

#include <sys/time.h>

double CpuSeconds() {
  struct timeval tp;
  gettimeofday(&tp, NULL); // 获取当前cpu时间戳。
  return ((double)tp.tv_sec + (double)tp.tv_usec * 1e-6);
}

参考：

man gettimeofday

2、检查cuda接口返回值

cuda的api会统一返回cudaError_t错误码，可以通过cudaGetErrorString()获取错误码对应的错误字符串。

将校验错误码的路径归纳为下面的宏，方便使用。

#define CHECK(call)       \
  do {                    \
    const cudaError_t error = call; \
    if (error != cudaSuccess) {     \
      printf("Error: %s: %d, ", __FILE__, __LINE__); \
      printf("code: %d, reason: %s\n", error, cudaGetErrorString(error)); \
      exit(1);            \
    }                     \
  } while(0)

参考：

3、cuda kernel实现

__global__ void kernelAdd(float *d_a, float *d_b, float *d_c, int n_elem) {
  int id = threadIdx.x + blockIdx.x * blockDim.x;

  // block * grid尺寸可能大于n_elem。超过n_elem的部分，就不计算了。
  if (id < n_elem) {
    d_c[id] = d_a[id] + d_b[id];
  }
}

4、调用cuda kernel

// 定义block和grid的尺寸，这里只使用了一维。
dim3 block(std::min(n_elem, max_thread_per_block));
dim3 grid((n_elem + block.x - 1) / block.x);

auto start_time = CpuSeconds();

// kernel调用是异步的，此句执行后会立即返回到host侧，需要手动调用cuda同步函数等待所有线程执行完毕。
kernelAdd<<<grid, block>>>(d_a, d_b, d_c, n_elem);

// 这一行非常重要！不做手动同步，测出的耗时只是kernel异步执行api返回到host的耗时，
// 并不是所有线程执行完的耗时。因此这里需要手动同步，等待所有线程执行结束。
cudaDeviceSynchronize(); 

printf("cuda add elapsed: %.6lfs\n", CpuSeconds() - start_time);

根据实际计算量划分block和grid的示意图如下：
grid_and_block

可以看到，虽然cuda生成了grid * block个线程，但最后一个block中的线程不一定会被全部用掉。

如何获取max_thread_per_block？

在host侧设置的gridDims和blockDims，受具体的硬件限制，在不同的GPU的上限是不同的。cuda提供了接口获取当前GPU硬件的上限信息。

代码如下：

cudaDeviceProp device_prop;
int dev = 0;
CHECK(cudaGetDeviceProperties(&device_prop, dev));
printf("Using device %d: %s, maxThreadsPerBlock: %d.\n\n",
       dev, device_prop.name, device_prop.maxThreadsPerBlock);
CHECK(cudaSetDevice(dev));

参考：

API手册：cudaGetDeviceProperties

5、使用nvprof测量耗时

nvprof加上所需执行的应用命令，即可执行profiling。本demo效果如下：

$ nvprof ./cuda_vector_add_timing
CUDA Demo: add two vectors.
==191428== NVPROF is profiling process 191428, command: ./cuda_vector_add_timing
Using device 0: NVIDIA GeForce GTX 1050 Ti, maxThreadsPerBlock: 1024.

testing n_elem: 16777216 ...
cpu add elapsed:  0.069825s
cuda block:(1024, 1, 1), grid:(16384, 1, 1)
cuda add elapsed: 0.002238s
cpu result : [1.735893, 0.446092, 1.168937, 1.654771, 1.396988, 1.141601, 0.817212, 1.561706, 0.467451, 1.386464, ...]
cuda result: [1.735893, 0.446092, 1.168937, 1.654771, 1.396988, 1.141601, 0.817212, 1.561706, 0.467451, 1.386464, ...]
cuda equals to cpu? yes.

==191428== Profiling application: ./cuda_vector_add_timing
==191428== Profiling result:
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
 GPU activities:   67.72%  39.002ms         1  39.002ms  39.002ms  39.002ms  [CUDA memcpy DtoH]
                   28.77%  16.571ms         2  8.2853ms  8.0925ms  8.4781ms  [CUDA memcpy HtoD]
                    3.51%  2.0214ms         1  2.0214ms  2.0214ms  2.0214ms  kernelAdd(float*, float*, float*, int)

计算两个16M数组的加法，CPU耗时70ms，GPU耗时2.23ms，是CPU的1/32。通过profiling结果可以看到，kernelAdd耗时2.02ms，比我们在host侧自己统计的耗时要短一点。这是因为nvprof能够严格测出device侧kernel执行耗时，而不必加上额外的host-device通信操作，更加精确。另外，host-device单次内存拷贝时间是8ms，是计算耗时的4倍。

参考：CUDA 专业提示：nvprof 是你便捷的通用 GPU 剖析器

小结

以上通过CPU时间戳和nvprof两种方式测量了kernel执行耗时。

本文示例完整代码请见：cuda_vector_add_timing.cu。

参考

[1] “Professional Cuda C Programming” Chapter 2.

本文采用署名-非商业性使用-相同方式共享 4.0 国际许可协议，转载请注明出处。