标签: nsight-compute

根据此链接https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html：

Warp 因等待 MIO（内存输入/输出）指令队列未满而停止。在 MIO 管道（包括特殊数学指令、动态分支以及共享内存指令）的极端利用率的情况下，这种停顿原因很严重。

根据此https://docs.nvidia.com/drive/drive_os_5.1.12.0L/nsight-graphics/activities/index.html：

可能由本地、全局、共享、属性、IPA、索引常量负载 (LDC) 和解耦数学触发。

我的理解是所有内存操作都是在LSU上执行的，所以我会想象它们一起存储在同一个指令队列中，然后由LSU单元执行。由于它们都一起排队，因此第二种解释（包括全局内存访问）对我来说更有意义。问题是，如果是这样的话，LG Throttle 就没有必要了。

MIO Throttle 实际上意味着什么？所有内存指令都存储在同一个队列中吗？

我正在学习__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 …

计算工作负载分析显示不同计算管道的利用率。我知道在现代 GPU 中，整数和浮点管道是不同的硬件单元，可以并行执行。但是，对于其他管道，哪个管道代表哪个硬件单元并不是很清楚。我也找不到任何关于管道缩写和解释的在线文档。

我的问题是：

1）ADU、CBU、TEX、XU的全称是什么？它们如何映射到硬件？

2) 哪些流水线使用相同的硬件单元（例如 FP16、FMA、FP64 使用浮点单元）？

3）现代GPU中的warp调度器每个周期可以调度2条指令（使用不同的管道）。哪些流水线可以同时使用（例如 FMA-ALU、FMA-SFU、ALU-Tensor 等）？

Ps：我为那些不熟悉Nsight Compute的人添加截图。

在最近的 NVIDIA 微架构中，有一种新的（？）warp 停顿原因/warp 调度程序状态分类法。

此分类法中的两个项目是：

• 短记分板- 记分板依赖于 MIO 队列操作。
• 长记分牌- 记分牌依赖于 L1TEX 操作。

我认为，“记分板”用于无序执行数据依赖性跟踪的意义（参见例如此处）。

我的问题：

• 形容词“短”或“长”描述什么？它是单个记分牌的长度吗？两种不同类型的操作的两个不同的记分牌？
• MIO 之间的这种有点不直观的二分法是什么意思 - 一些，但不是全部都是内存操作；和 L1TEX 操作，哪些都是内存操作？这是一个二分法wrt停顿原因还是与实际硬件有关？

两个问题：

1. 根据 Nsight Compute 的说法，我的内核是受计算限制的。相对于峰值性能的 SM 利用率为 74%，内存利用率为 47%。然而，当我查看每个管道利用率时，LSU 利用率远高于其他管道（75% 与 10-15%）。这难道不表明我的内核受内存限制吗？如果计算和内存资源的利用率与管道利用率不对应，我不知道如何解释这些术语。

2. 调度程序每 4 个周期才发出一次，这是否意味着我的内核受到延迟限制？人们通常根据计算和内存资源的利用率来定义它。两者之间是什么关系？

假设我有一个myapp不需要命令行参数的可执行文件，并启动 CUDA 内核mykernel。我可以调用：

nv-nsight-cu-cli -k mykernel myapp

并得到如下所示的输出：

==PROF== Connected to process 30446 (/path/to/myapp)
==PROF== Profiling "mykernel": 0%....50%....100% - 13 passes
==PROF== Disconnected from process 1234
[1234] myapp@127.0.0.1
  mykernel(), 2020-Oct-25 01:23:45, Context 1, Stream 7
    Section: GPU Speed Of Light
    --------------------------------------------------------------------
    Memory Frequency                      cycle/nsecond      1.62
    SOL FB                                %                  1.58
    Elapsed Cycles                        cycle              4,421,067
    SM Frequency                          cycle/nsecond      1.43
    Memory [%]                            %                  61.76
    Duration                              msecond            3.07
    SOL L2                                %                  0.79
    SM Active Cycles                      cycle              4,390,420.69
    (etc. etc.)
    -------------------------------------------------------------------- …