在 GPU 执行之外启动 CUDA 内核的大量开销

Question

我通过测量从启动内核之前到 cudaDeviceSynchronize 之后的间隔（使用 gettimeofday）来测量内核的运行时间，从 CPU 线程来看。在开始记录间隔之前，我有一个 cudaDeviceSynchronize。我还检测内核以在内核开始时通过线程（0,0,0）从块（0,0,0）到块（占用1， 0,0) 到一个大小等于 SM 数量的数组。内核代码末尾的每个线程都会将时间戳更新到另一个数组（大小相同），其索引等于它运行的 SM 的索引。

从两个数组计算的间隔是从 CPU 线程测量的间隔的 60-70%。

例如，在 K40 上，虽然 gettimeofday 给出的间隔为 140 毫秒，但从 GPU 时间戳计算的间隔平均值仅为 100 毫秒。我已经尝试了许多网格大小（15 块到 6K 块），但到目前为止发现了类似的行为。

__global__ void some_kernel(long long *d_start, long long *d_end){
     if(threadIdx.x==0){
        d_start[blockIdx.x] = clock64();
     }
     //some_kernel code
     d_end[blockIdx.x] = clock64();
}

这在专家看来可行吗？

Answer 1

这在专家看来可行吗？

我想对于您未显示的代码，一切皆有可能。毕竟，您的任何计算算法中都可能有一个愚蠢的错误。但是，如果问题是“对于需要约 140 毫秒执行的内核，在内核启动时应该有 40 毫秒的未计入时间开销是否明智？”我会说不。

我相信我在 cmets 中概述的方法相当准确。从网格中的任何线程获取最小clock64() 时间戳（但请参阅下面有关 SM 限制的注释）。将其与网格中任何线程的最大时间戳进行比较。根据我的测试，差异将与报告的 gettimeofday() 执行时间相差 2% 以内。

这是我的测试用例：

$ cat t1040.cu
#include <stdio.h>
#include <stdlib.h>
#include <stdint.h>

#define LS_MAX 2000000000U
#define MAX_SM 64

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)


#include <time.h>
#include <sys/time.h>
#define USECPSEC 1000000ULL

__device__ int result;
__device__ unsigned long long t_start[MAX_SM];
__device__ unsigned long long t_end[MAX_SM];

unsigned long long dtime_usec(unsigned long long start){

  timeval tv;
  gettimeofday(&tv, 0);
  return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start;
}

__device__ __inline__ uint32_t __mysmid(){
  uint32_t smid;
  asm volatile("mov.u32 %0, %%smid;" : "=r"(smid));
  return smid;}

__global__ void kernel(unsigned ls){

  unsigned long long int ts = clock64();
  unsigned my_sm = __mysmid();
  atomicMin(t_start+my_sm, ts);
  // junk code to waste time
  int tv = ts&0x1F;
  for (unsigned i = 0; i < ls; i++){
    tv &= (ts+i);}
  result = tv;
  // end of junk code
  ts = clock64();
  atomicMax(t_end+my_sm, ts);

}

// optional command line parameter 1 = kernel duration, parameter 2 = number of blocks, parameter 3 = number of threads per block
int main(int argc, char *argv[]){

 unsigned ls;
 if (argc > 1) ls = atoi(argv[1]);
 else ls = 1000000;
 if (ls > LS_MAX) ls = LS_MAX;
 int num_sms = 0;
 cudaDeviceGetAttribute(&num_sms, cudaDevAttrMultiProcessorCount, 0);
 cudaCheckErrors("cuda get attribute fail");
 int gpu_clk = 0;
 cudaDeviceGetAttribute(&gpu_clk, cudaDevAttrClockRate, 0);
 if ((num_sms < 1) || (num_sms > MAX_SM)) {printf("invalid sm count: %d\n", num_sms); return 1;}
 unsigned blks;
 if (argc > 2) blks = atoi(argv[2]);
 else blks = num_sms;
 if ((blks < 1) || (blks > 0x3FFFFFFF)) {printf("invalid blocks: %d\n", blks); return 1;}
 unsigned ntpb;
 if (argc > 3) ntpb = atoi(argv[3]);
 else ntpb = 256;
 if ((ntpb < 1) || (ntpb > 1024)) {printf("invalid threads: %d\n", ntpb); return 1;}
 kernel<<<1,1>>>(100);  // warm up
 cudaDeviceSynchronize();
 cudaCheckErrors("kernel fail");
 unsigned long long *h_start, *h_end;
 h_start = new unsigned long long[num_sms];
 h_end = new unsigned long long[num_sms];
 for (int i = 0; i < num_sms; i++){
   h_start[i] = 0xFFFFFFFFFFFFFFFFULL;
   h_end[i] = 0;}
 cudaMemcpyToSymbol(t_start, h_start, num_sms*sizeof(unsigned long long));
 cudaMemcpyToSymbol(t_end, h_end, num_sms*sizeof(unsigned long long));
 unsigned long long htime = dtime_usec(0);
 kernel<<<blks,ntpb>>>(ls);
 cudaDeviceSynchronize();
 htime = dtime_usec(htime);
 cudaMemcpyFromSymbol(h_start, t_start, num_sms*sizeof(unsigned long long));
 cudaMemcpyFromSymbol(h_end, t_end, num_sms*sizeof(unsigned long long));
 cudaCheckErrors("some error");
 printf("host elapsed time (ms): %f \n device sm clocks:\n start:", htime/1000.0f);
 unsigned long long max_diff = 0;
 for (int i = 0; i < num_sms; i++) {printf(" %12lu  ", h_start[i]);}
 printf("\n end:  ");
 for (int i = 0; i < num_sms; i++) {printf(" %12lu  ", h_end[i]);}
 for (int i = 0; i < num_sms; i++) if ((h_start[i] != 0xFFFFFFFFFFFFFFFFULL) && (h_end[i] != 0) && ((h_end[i]-h_start[i]) > max_diff)) max_diff=(h_end[i]-h_start[i]);
 printf("\n max diff clks: %lu\nmax diff kernel time (ms): %f\n", max_diff, max_diff/(float)(gpu_clk));
 return 0;
}
$ nvcc -o t1040 t1040.cu -arch=sm_35
$ ./t1040 1000000 1000 128
host elapsed time (ms): 2128.818115
 device sm clocks:
 start:      3484744        3484724
 end:     2219687393     2228431323
 max diff clks: 2224946599
max diff kernel time (ms): 2128.117432
$

注意事项：

1. 由于使用 64 位 atomicMin 和 atomicMax，此代码只能在 cc3.5 或更高版本的 GPU 上运行。

2. 我已经在 GT640（非常低端 cc3.5 设备）和 K40c（高端）上运行了各种网格配置，主机和设备之间的时序结果一致在 2% 以内（对于相当长的内核执行时间。如果您将1 作为命令行参数传递，网格尺寸非常小，内核执行时间将非常短（纳秒），而主机将看到大约 10-20us。这 is 正在测量内核启动开销。因此 2% 的数字是用于执行时间超过 20us 的内核）。

3. 它接受 3 个（可选）命令行参数，第一个参数会改变内核执行的时间。

4. 我的时间戳是基于 每个 SM 完成的，因为 clock64() 资源是 indicated to be 一个 每个 SM 资源。 SM 时钟不保证在 SM 之间同步。

5. 您可以修改网格尺寸。第二个可选命令行参数指定要启动的块数。第三个可选命令行参数指定每个块的线程数。我在这里展示的计时方法不应依赖于启动的块数或每个块的线程数。如果您指定的块少于 SM，则代码应忽略“未使用”的 SM 数据。