cuda - Nsight Compute 如何确定/显示共享内存指标？

Question

我正在学习 __shared__ CUDA 中的内存，我对 Nsight Compute 如何显示共享内存统计信息感到困惑。

我正在经历this article (代码在 Nvidia 的 github here 上可用，但在下面复制以供引用)。

#include <stdio.h>

__global__ void staticReverse(int *d, int n)
{
  __shared__ int s[64];
  int t = threadIdx.x;
  int tr = n-t-1;
  s[t] = d[t];
  __syncthreads();
  d[t] = s[tr];
}

__global__ void dynamicReverse(int *d, int n)
{
  extern __shared__ int s[];
  int t = threadIdx.x;
  int tr = n-t-1;
  s[t] = d[t];
  __syncthreads();
  d[t] = s[tr];
}

int main(void)
{
  const int n = 64;
  int a[n], r[n], d[n];

  for (int i = 0; i < n; i++) {
    a[i] = i;
    r[i] = n-i-1;
    d[i] = 0;
  }

  int *d_d;
  cudaMalloc(&d_d, n * sizeof(int)); 

  // run version with static shared memory
  cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
  staticReverse<<<1,n>>>(d_d, n);
  cudaMemcpy(d, d_d, n*sizeof(int), cudaMemcpyDeviceToHost);
  for (int i = 0; i < n; i++) 
    if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);

  // run dynamic shared memory version
  cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
  dynamicReverse<<<1,n,n*sizeof(int)>>>(d_d, n);
  cudaMemcpy(d, d_d, n * sizeof(int), cudaMemcpyDeviceToHost);
  for (int i = 0; i < n; i++) 
    if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);
}

当我运行 Nsight Compute 时，我看到以下图表 staticReverse内核( dynamicReverse 内核几乎相同):



问题 1 : 图表显示1个请求到共享内存和1个请求来自共享内存，但为什么它也显示0个共享内存指令？请求不算作指令吗？从这个图表的角度来看，什么算作共享内存指令？

接下来，在源 View 中，Nsight Compute 显示各种指标的逐行计数:



问题 2 :为什么第 8 行和第 10 行的“Memory L1 Transactions Shared”显示为 0？我期待看到:

第 8 行:相等数量的 [从全局内存加载事务] 和 [将事务存储到共享内存]

第 10 行:相等数量的 [从共享内存加载事务] 和 [将事务存储到全局内存]

问题 3 :为什么第 8 行和第 10 行各有 8 个内存事务？

我的系统:

Ubuntu 18.04 LTS

GeForce 1070 (帕斯卡)

CUDA 版本:10.2

驱动程序版本:440.64.00

Answer 1

如果您可以检查(并在此处显示)Source 页面的低级 SASS View 以及高级 CUDA-C View ，那就太好了。根据 SASS(汇编)指令收集源指标，然后在 CUDA-C View 中汇总。检查实际程序集可以提供有关编译器生成的指令类型的信息，并且可以更好地解释您看到的数据。

Does a request not count as an instruction? From this chart's perspective, what counts as a shared memory instruction?

请求和指示不是一回事。指令是正在执行的实际 SASS 汇编指令。请求由硬件作为执行指令的结果生成，请求的数量可能会根据代码的表现而有所不同。