2. nsights 可视化分析GPU核函数性能

Bertha 。 2022-12-23 03:57 289阅读 0赞

例子：

#include <stdio.h>
    
    void initWith(float num, float *a, int N)
    {
      for(int i = 0; i < N; ++i)
      {
        a[i] = num;
      }
    }
    
    __global__
    void addVectorsInto(float *result, float *a, float *b, int N)
    {
      int index = threadIdx.x + blockIdx.x * blockDim.x;
      int stride = blockDim.x * gridDim.x;
    
      for(int i = index; i < N; i += stride)
      {
        result[i] = a[i] + b[i];
      }
    }
    
    void checkElementsAre(float target, float *vector, int N)
    {
      for(int i = 0; i < N; i++)
      {
        if(vector[i] != target)
        {
          printf("FAIL: vector[%d] - %0.0f does not equal %0.0f\n", i, vector[i], target);
          exit(1);
        }
      }
      printf("Success! All values calculated correctly.\n");
    }
    
    int main()
    {
      int deviceId;
      int numberOfSMs;
    
      cudaGetDevice(&deviceId);
      cudaDeviceGetAttribute(&numberOfSMs, cudaDevAttrMultiProcessorCount, deviceId);
    
      const int N = 2<<24;
      size_t size = N * sizeof(float);
    
      float *a;
      float *b;
      float *c;
    
      cudaMallocManaged(&a, size);
      cudaMallocManaged(&b, size);
      cudaMallocManaged(&c, size);
    
      initWith(3, a, N);
      initWith(4, b, N);
      initWith(0, c, N);
    
      size_t threadsPerBlock;
      size_t numberOfBlocks;
    
      threadsPerBlock = 256;
      numberOfBlocks = 32 * numberOfSMs;
    
      cudaError_t addVectorsErr;
      cudaError_t asyncErr;
    
      addVectorsInto<<<numberOfBlocks, threadsPerBlock>>>(c, a, b, N);
    
      addVectorsErr = cudaGetLastError();
      if(addVectorsErr != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(addVectorsErr));
    
      asyncErr = cudaDeviceSynchronize();
      if(asyncErr != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(asyncErr));
    
      checkElementsAre(7, c, N);
    
      cudaFree(a);
      cudaFree(b);
      cudaFree(c);
    }

编译：

nvcc -o vector-add-no-prefetch 01-vector-add/01-vector-add.cu -run

生成报告：

nsys profile --stats=true -o vector-add-no-prefetch-report ./vector-add-no-prefetch

使用nsight-sys 打开文件  
![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70_pic_center]  
通过命令行打开nsight-sys ，打开刚刚生成的qdrep文件  
![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70_pic_center 1]  
主要查看CUDA部分  
![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70_pic_center 2]

![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70_pic_center 3]

[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70_pic_center]: /images/20221120/d340cbba384f481699c3df80457216dc.png
[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70_pic_center 1]: https://img-blog.csdnimg.cn/20201122154639454.png?x-oss-process=image/watermark,type_ZmFuZ3poZW5naGVpdGk,shadow_10,text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ==,size_16,color_FFFFFF,t_70#pic_center
[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70_pic_center 2]: https://img-blog.csdnimg.cn/20201122155338600.png?x-oss-process=image/watermark,type_ZmFuZ3poZW5naGVpdGk,shadow_10,text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ==,size_16,color_FFFFFF,t_70#pic_center
[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70_pic_center 3]: https://img-blog.csdnimg.cn/2020112215520149.png?x-oss-process=image/watermark,type_ZmFuZ3poZW5naGVpdGk,shadow_10,text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ==,size_16,color_FFFFFF,t_70#pic_center