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

Posted 2023-03-06

【中文标题】Nsight Compute 如何确定/显示共享内存指标？

【英文标题】：How does Nsight Compute determine/display shared memory metrics?

【发布时间】：2020-05-27 11:55:49

【问题描述】：

我正在了解 CUDA 中的 __shared__ 内存，但我对 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 个共享内存指令？请求不算作指令吗？从这个图表的角度来看，什么算作共享内存指令？

接下来，在源代码视图中，Nsight Compute 显示各种指标的逐行计数：

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

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

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

我的系统：

Ubuntu 18.04 LTS GeForce 1070（帕斯卡） CUDA 版本：10.2 驱动版本：440.64.00

【问题讨论】：

我打算研究 PTX 文件并将它们链接到报告的两个内核的 Nsight 性能指标，以试图回答您的问题，但意识到我实际上无法重现与您相同的内存图表.就我而言，我看到的共享内存 instructions 数量与两个内核的全局内存相同。我有 CUDA 10.2 并针对 CC 6.1 但在 Windows 10 上编译了代码。

【参考方案1】：

如果您可以检查（并在此处显示）源页面的低级 SASS 视图以及高级 CUDA-C 视图，那就太好了。源指标是按 SASS（汇编）指令收集的，然后在 CUDA-C 视图中汇总。检查实际程序集可以了解编译器生成的指令类型，并且可以更好地解释您所看到的数据。

请求不算作指令吗？从这个图表的角度来看，什么算作共享内存指令？

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

【讨论】：

我看看能不能在今天晚些时候发布 SASS。回复：指令与请求，我知道 1 条指令如何生成 1 个或更多请求，但我不清楚 0 条指令中怎么可能有 1 个请求。