两个相同的空内核上的奇怪 cuda 性能。答案

【问题标题】：Weird cuda performance on two same empty kernels.两个相同的空内核上的奇怪 cuda 性能。
【发布时间】：2018-10-21 15:39:25
【问题描述】：

我有两个“空”内核，每个内核都有一个永远不会被触及的 if 语句。

#include <cstdio>
#include <time.h>
#include <sys/time.h>
#include <cuda.h>

inline double wtime(){
  double time[2];
  struct timeval time1;
  gettimeofday(&time1, NULL);

  time[0]=time1.tv_sec;
  time[1]=time1.tv_usec;

  return time[0]+time[1]*1.0e-6;
}

__global__ void __empty1(bool flag){if(flag){ printf("hh\n");}}
__global__ void __empty2(bool flag){if(flag){ ; }}

int main(){

  cudaDeviceSynchronize();
  double s = wtime();
  __empty1<<<256,256>>>(false);
  cudaDeviceSynchronize();
  printf("empty1: %.3f\n", 1000*(wtime()-s));

  cudaDeviceSynchronize();
  s = wtime();
  __empty2<<<256,256>>>(false);
  cudaDeviceSynchronize();
  printf("empty2: %.3f\n", 1000*(wtime()-s));
  return 0;
}

我用cuda-7.5和-O3编译了代码，然后在K40m上运行。

nvcc -O3 -arch=sm_35 ./main.cu

第一个空内核需要 1ms，而第二个内核需要 0.02ms。

empty1: 1.075
empty2: 0.019

这很奇怪，因为这两个内核都没有进入分支。这两个内核的运行时间预计是相同的。

【问题讨论】：

是什么让你认为运行时间是一样的？可能有很多无条件包含和执行的 printf 支持代码。除非您已经反汇编并检查了 GPU 在这两种情况下运行的实际 SASS 代码，否则您无法做出这样的假设

标签： performance cuda

【解决方案1】：

代码不一样。正如 cmets 中所指出的，这可以通过查看 SASS 来发现。

例子：

$ cat t1353.cu
#include <cstdio>
#include <time.h>
#include <sys/time.h>
#include <cuda.h>

inline double wtime(){
  double time[2];
  struct timeval time1;
  gettimeofday(&time1, NULL);

  time[0]=time1.tv_sec;
  time[1]=time1.tv_usec;

  return time[0]+time[1]*1.0e-6;
}

__global__ void __empty1(bool flag){if(flag){ printf("hh\n");}}
__global__ void __empty2(bool flag){if(flag){ ; }}

int main(){

  cudaDeviceSynchronize();
  double s = wtime();
  __empty1<<<256,256>>>(false);
  cudaDeviceSynchronize();
  printf("empty1: %.3f\n", 1000*(wtime()-s));

  cudaDeviceSynchronize();
  s = wtime();
  __empty2<<<256,256>>>(false);
  cudaDeviceSynchronize();
  printf("empty2: %.3f\n", 1000*(wtime()-s));
  return 0;
}
$ nvcc -arch=sm_35 -o t1353 t1353.cu
$ cuobjdump -sass t1353

Fatbin elf code:
================
arch = sm_35
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit

        code for sm_35

Fatbin elf code:
================
arch = sm_35
code version = [1,7]
producer = cuda
host = linux
compile_size = 64bit

        code for sm_35
                Function : _Z8__empty2b
        .headerflags    @"EF_CUDA_SM35 EF_CUDA_PTX_SM(EF_CUDA_SM35)"
                                                          /* 0x0800000000b81000 */
        /*0008*/                   MOV R1, c[0x0][0x44];  /* 0x64c03c00089c0006 */
        /*0010*/                   MOV RZ, RZ;            /* 0xe4c03c007f9c03fe */
        /*0018*/                   EXIT;                  /* 0x18000000001c003c */
        /*0020*/                   BRA 0x20;              /* 0x12007ffffc1c003c */
        /*0028*/                   NOP;                   /* 0x85800000001c3c02 */
        /*0030*/                   NOP;                   /* 0x85800000001c3c02 */
        /*0038*/                   NOP;                   /* 0x85800000001c3c02 */
                .............................


                Function : _Z8__empty1b
        .headerflags    @"EF_CUDA_SM35 EF_CUDA_PTX_SM(EF_CUDA_SM35)"
                                                                     /* 0x08b8b0a0a0a0a000 */
        /*0008*/                   MOV R1, c[0x0][0x44];             /* 0x64c03c00089c0006 */
        /*0010*/                   LDC.S8 R0, c[0x0][0x140];         /* 0x7c880000a01ffc02 */
        /*0018*/                   I2I.U16.S8 R0, R0;                /* 0xe6000000001c8402 */
        /*0020*/                   LOP.AND R0, R0, 0xff;             /* 0xc20000007f9c0001 */
        /*0028*/                   I2I.S32.S16 R0, R0;               /* 0xe6000000001cd802 */
        /*0030*/                   ISETP.EQ.AND P0, PT, R0, RZ, PT;  /* 0xdb281c007f9c001e */
        /*0038*/               @P0 EXIT;                             /* 0x180000000000003c */
                                                                     /* 0x08b810b800108010 */
        /*0048*/                   MOV32I R4, 0x0;                   /* 0x74000000001fc012 */
        /*0050*/                   MOV32I R5, 0x0;                   /* 0x74000000001fc016 */
        /*0058*/                   MOV R7, RZ;                       /* 0xe4c03c007f9c001e */
        /*0060*/                   MOV R6, RZ;                       /* 0xe4c03c007f9c001a */
        /*0068*/                   JCAL 0x0;                         /* 0x1100000000000100 */
        /*0070*/                   MOV RZ, RZ;                       /* 0xe4c03c007f9c03fe */
        /*0078*/                   EXIT;                             /* 0x18000000001c003c */
        /*0080*/                   BRA 0x80;                         /* 0x12007ffffc1c003c */
        /*0088*/                   NOP;                              /* 0x85800000001c3c02 */
        /*0090*/                   NOP;                              /* 0x85800000001c3c02 */
        /*0098*/                   NOP;                              /* 0x85800000001c3c02 */
        /*00a0*/                   NOP;                              /* 0x85800000001c3c02 */
        /*00a8*/                   NOP;                              /* 0x85800000001c3c02 */
        /*00b0*/                   NOP;                              /* 0x85800000001c3c02 */
        /*00b8*/                   NOP;                              /* 0x85800000001c3c02 */
                .............................



Fatbin ptx code:
================
arch = sm_35
code version = [5,0]
producer = cuda
host = linux
compile_size = 64bit
compressed
$

所以在empty2 的情况下，机器代码如下所示：

        /*0008*/                   MOV R1, c[0x0][0x44];  /* 0x64c03c00089c0006 */
        /*0010*/                   MOV RZ, RZ;            /* 0xe4c03c007f9c03fe */
        /*0018*/                   EXIT;                  /*

empty1 情况下更长：

        /*0008*/                   MOV R1, c[0x0][0x44];             /* 0x64c03c00089c0006 */
        /*0010*/                   LDC.S8 R0, c[0x0][0x140];         /* 0x7c880000a01ffc02 */
        /*0018*/                   I2I.U16.S8 R0, R0;                /* 0xe6000000001c8402 */
        /*0020*/                   LOP.AND R0, R0, 0xff;             /* 0xc20000007f9c0001 */
        /*0028*/                   I2I.S32.S16 R0, R0;               /* 0xe6000000001cd802 */
        /*0030*/                   ISETP.EQ.AND P0, PT, R0, RZ, PT;  /* 0xdb281c007f9c001e */
        /*0038*/               @P0 EXIT;                            

...
        /*0078*/                   EXIT;                             /* 0x18000000001c003c */

这里更大的问题可能是时间严格性/正确性之一。 CUDA 有延迟初始化。这意味着您的 CUDA 代码中的第一组调用可能会产生比通常数量更多的时间开销。根据我的测试，这会影响这里的结论。如果我在实际计时之前对empty1 进行“热身”调用，则两种情况之间的测量时间几乎相同。这可以通过代码长度的差异来合理解释。

例子：

$ cat t1353.cu
#include <cstdio>
#include <time.h>
#include <sys/time.h>
#include <cuda.h>

inline double wtime(){
  double time[2];
  struct timeval time1;
  gettimeofday(&time1, NULL);

  time[0]=time1.tv_sec;
  time[1]=time1.tv_usec;

  return time[0]+time[1]*1.0e-6;
}

__global__ void __empty1(bool flag){if(flag){ printf("hh\n");}}
__global__ void __empty2(bool flag){if(flag){ ; }}

int main(){

  __empty1<<<256,256>>>(false);
  cudaDeviceSynchronize();
  double s = wtime();
  __empty1<<<256,256>>>(false);
  cudaDeviceSynchronize();
  printf("empty1: %.3f\n", 1000*(wtime()-s));

  __empty2<<<256,256>>>(false);
  cudaDeviceSynchronize();
  s = wtime();
  __empty2<<<256,256>>>(false);
  cudaDeviceSynchronize();
  printf("empty2: %.3f\n", 1000*(wtime()-s));
  return 0;
}
$ nvcc -arch=sm_35 -o t1353 t1353.cu
$ ./t1353
empty1: 0.023
empty2: 0.015
$

【讨论】：