我正在学习__shared__CUDA 中的内存，但我对 Nsight Compute 如何显示共享内存统计信息感到困惑。
我正在阅读这篇文章（代码在 Nvidia 的 github 上可用，但在下面复制以供参考）。
#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
