memory - NVPROF 报告的交易指标到底是什么？

Question

我试图弄清楚“nvprof”报告的每个指标到底是什么。更具体地说，我无法弄清楚哪些事务是系统内存和设备内存的读写。我编写了一个非常基本的代码只是为了帮助解决这个问题。

#define TYPE float
#define BDIMX 16
#define BDIMY 16
#include <cuda.h>
#include <cstdio>
#include <iostream>
__global__ void kernel(TYPE *g_output, TYPE *g_input, const int dimx, const int dimy)
{
__shared__ float s_data[BDIMY][BDIMX];
  int ix = blockIdx.x * blockDim.x + threadIdx.x;
  int iy = blockIdx.y * blockDim.y + threadIdx.y;
  int in_idx = iy * dimx + ix; // index for reading input
  int tx = threadIdx.x; // thread’s x-index into corresponding shared memory tile  
  int ty = threadIdx.y; // thread’s y-index into corresponding shared memory tile 
  s_data[ty][tx] = g_input[in_idx];
  __syncthreads();
  g_output[in_idx] = s_data[ty][tx] * 1.3;
  }


int main(){
  int size_x = 16, size_y = 16;
  dim3 numTB;
    numTB.x = (int)ceil((double)(size_x)/(double)BDIMX) ;
    numTB.y = (int)ceil((double)(size_y)/(double)BDIMY) ;
  dim3 tbSize; 
  tbSize.x = BDIMX;
  tbSize.y = BDIMY;
  float* a,* a_out;
  float *a_d = (float *) malloc(size_x * size_y * sizeof(TYPE));
  cudaMalloc((void**)&a,     size_x * size_y * sizeof(TYPE));
  cudaMalloc((void**)&a_out, size_x * size_y * sizeof(TYPE));
  for(int index = 0; index < size_x * size_y; index++){
      a_d[index] = index;
   }
  cudaMemcpy(a, a_d, size_x * size_y * sizeof(TYPE), cudaMemcpyHostToDevice);
  kernel <<<numTB, tbSize>>>(a_out, a, size_x, size_y);
  cudaDeviceSynchronize();
  return 0;
}

然后我运行 nvprof --metrics all 以获得输出以查看所有指标。这是我感兴趣的部分:

           Metric Name                        Metric Description         Min         Max         Avg
Device "Tesla K40c (0)"
  Kernel: kernel(float*, float*, int, int)
    local_load_transactions                   Local Load Transactions           0           0           0
   local_store_transactions                  Local Store Transactions           0           0           0
   shared_load_transactions                  Shared Load Transactions           8           8           8
  shared_store_transactions                 Shared Store Transactions           8           8           8
           gld_transactions                  Global Load Transactions           8           8           8
           gst_transactions                 Global Store Transactions           8           8           8
   sysmem_read_transactions           System Memory Read Transactions           0           0           0
  sysmem_write_transactions          System Memory Write Transactions           4           4           4
     tex_cache_transactions                Texture Cache Transactions           0           0           0
     dram_read_transactions           Device Memory Read Transactions           0           0           0
    dram_write_transactions          Device Memory Write Transactions          40          40          40
       l2_read_transactions                      L2 Read Transactions          70          70          70
      l2_write_transactions                     L2 Write Transactions          46          46          46

我了解共享和全局访问。全局访问被合并，并且由于有 8 个扭曲，因此有 8 个事务。但我无法弄清楚系统内存和设备内存写入事务号。

Answer 1

如果您拥有一个同时具有逻辑和物理空间的 GPU 内存层次结构模型，例如 here，这会很有帮助。 .

引用“概览选项卡”图:

1. gld_transactions 是指从 warp 发出的针对全局逻辑空间的事务。在图中，这将是从左侧的“内核”框到右侧的“全局”框的线，逻辑数据移动方向将从右到左。

2. gst_transactions 引用与上面相同的行，但逻辑上是从左到右。请注意，这些逻辑全局事务可能会命中缓存，然后不会再去任何地方。从指标的角度来看，这些交易类型仅指图表上指示的线。

3. dram_write_transactions 是指图中右侧设备内存与二级缓存连接的线，该线上的逻辑数据流向是从左到右。由于 L2 缓存行为 32 字节(而 L1 缓存行和全局事务的大小为 128 字节)，因此设备内存事务也是 32 字节，而不是 128 字节。因此，经过 L1(如果启用，它是直写式缓存)和 L2 的全局写事务将生成 4 个 dram_write 事务。这应该可以解释 40 笔交易中的 32 笔。

4. 系统内存事务以零拷贝主机内存为目标。你似乎没有这个，所以我无法解释这些。

请注意，在某些情况下，对于某些指标，在某些 GPU 上，分析器在启动非常少量的线程 block 时可能会出现一些“不准确”。例如，某些指标是在每个 SM 的基础上进行采样并缩放的。 (但是，设备内存事务不属于此类)。如果您在每个 SM 上完成不同的工作(可能是由于启动的线程 block 数量非常少)，那么缩放可能会产生误导/不太准确。一般来说，如果您启动大量线程 block ，这些线程 block 通常会变得无关紧要。

This answer可能也会感兴趣。