Nsight Graphics

之前分析GPU系统的时候，用的多是Nsight System。但目前发现Nsight Graphics可能更适合分析一段时间内的渲染细节，以及对应的GPU行为。苦于看不懂界面，于是决定把一些学习的记录整理下来。

1. Self Instruction Mix

1. 基本执行单元

首先，从认识基本的执行单元开始，以下是一览表：

缩写	含义	实际作用
FMA	Fused Multiply-Add/Accumulate	主要执行 FP32 加、乘、融合乘加，也执行部分整数乘法/乘加。
ALU	Arithmetic Logic Unit	普通整数运算、比较、逻辑、位操作、数据移动，以及部分简单 FP32 操作。
LDCU	常量加载到 Uniform Register 的指令/通路分类	从 Constant Memory 向 warp-uniform 寄存器加载标量值。官方对它的定义是“Load a Value from Constant Memory into a Uniform Register”。CUDA 指令说明
UDP	Uniform Data Path	统一/标量数据通路。处理整个 warp 都相同的值，避免每个 lane 重复做同一计算。新版指标中也常称 `UNIFORM`。
ADU	Address Divergence Unit	地址分歧处理、部分跳转/分支、常量加载、CTA/block barrier 等控制和寻址工作。
SCH	Scheduler / 调度控制路径	调度和控制类微操作、NOP，以及部分光追控制指令。它不等于 RT Core 本体。
CBU	Convergence Barrier Unit	warp 分支汇合、控制流、barrier、branch/call/return/exit 等。
LSU	Load/Store Unit	发出 Global、Local、Shared Memory 的 load/store/atomic，以及部分特殊寄存器、warp primitive 和 barrier 指令。
XU	Transcendental and Data Type Conversion Unit	特殊函数和类型转换，如 `sin/cos/rcp/rsqrt`、float↔int；经常也被称为 SFU 路径。
TEX	Texture Unit	纹理寻址、LOD、过滤、格式转换、Texture/Surface 访问。
RTCORE	Ray Tracing Core	BVH 遍历、光线与包围盒/三角形求交等硬件光追工作；不是整个 ray tracing shader 的全部执行。
IPA	Pixel Shader 属性插值/访问路径	读取或插值 Pixel Shader 输入属性。公开文档一般直接使用 IPA，没有稳定公开的全称。不要与编译器中的 Interprocedural Analysis 混淆。
REDUX	Reduction 路径/指令	对 warp 中多个 lane 的值执行归约，并产生 uniform 结果。

接下来展开一些细节：

FMA & ALU

Q：为什么FMA要把乘法和加法统合在一起？不能先乘再加吗？

A：先乘再加，中间结果和最后结果一共要rounded 2次；FMA 只 rounded 一次，数值精度更高

但在算FLOPS（Floating point operations per second）时，FMA算两次浮点运算

浮点乘法加法实际上复用了FMA的结构，可以理解为 x * y + 0这样的形式（当然不一定真是这样编码）

Q：FMA 和 ALU的职责范围？为什么一定要分离？

核心：乘法比加法/逻辑/比较贵的多

从这两张实际Nsight截下来的图可以大致看出两个pipe的功能（以及在一帧渲染中的大概指令数）

记录一些个人的疑点：

Q：Logic Operation 和 FP32 Logic的区别？

Logic Op：按位与，或，非，异或，bool运算等，每一位独立参加运算，不在意语义
FP32 Logic：min/max、选择（float r2 = condition ? a : b;）、saturate（clamp），需要考虑语义

Q：Compare 和 Logic的区别？

Compare：小于大于等于，输出bool值，代表两个数的相对关系，因此没有第三个数
Logic Op：与或非，不关心输出什么，可以输出bool值（如&不输出bool值，而&&输出bool值），关心结果数

Q：什么是Data Movement？

赋值（a = b），三元选择（），swizzle，打包

Q：FMA int乘加和 ALU int计算的边界？

A：所有需要乘法的int 计算用FMA；其他的所有都用ALU

Q：Data Movement是什么？

A：类似指令集里的MOV，表示寄存器 to 寄存器，或者立即数 to 寄存器

LSU

LSU代表着显存和寄存器的交互，主要分Global Shared Local三级

这里有一个误区：我们都知道GPU有三级缓存，L2（一块GPU的所有SM共享），L1（一个SM的所有subcore共享），L0（一个subcore上的所有FP 32/INT 32和一个Tensor core共享），但他们和LSU这里的三级不是一一对应的。

LSU这里对应的，更像是shader可见的虚拟地址空间，其与实际物理空间的对应是不确定的。Cache仅仅是加速结构，比如实际上是不能直接分配L2 Cache的，只能在Load VM的时候加载到L2 Cache

Global Memory是最常见的memory，基本所有没有特殊设置的显存操作，都是对Global memory的操作，一般来说，对于Global Memory的instruction和Sample均远高于Shared 和 Local

Shared Memory很特殊，虽然之前说过Cache是不能分配的，但Shared Memory确实和L1 Cache共享同样的物理空间，属于所谓的 On-chip 片上内存，Load/Store 都相当的快。因此常用于特殊的shader优化，比如一个warp共享的数据可以放在shared memory中（TODO：check this）

Local Memory 是寄存器资源耗尽后的备选方案：线程的私有数据本来应该全部放在寄存器中，但如果寄存器资源耗尽，就会被放在显存中。一般用来观察寄存器资源是否比较极限。

GPU架构：以Blackwell架构为例，一个subcore含有16个FP32专用lane，以及16个可选FP32/INT32 lane，宣传的时候称为32个 cuda core。一个SM有4个subcore，因此有128个cuda core。

一般来说，对于Global Memory的调用次数和Sample均远高于Shared 和 Local

LDCU

Load a Value from Constant Memory into a Uniform Register

Q：什么是Constant Memory？Constant Memory 和 Shared Memory 的区别？

这两种 Memory 理论上都是所谓片上缓存（on-chip就是SM内部）。

shared memory

L1 cache的一部分，是真正的片上缓存
仅仅在一个block中的所有线程共享
可读可写
Shared memory是SM的稀缺资源，使用过多会限制SM上同时驻留的线程块数量（Occupancy）
以RTX PRO 6000为例，一个SM上面有 256 KB的L1 Cache/Shared memory共用区域，在不做特殊调整的情况下，L1 Cache 和 Shared memory是各128 KB，但可以通过函数来设置偏好L1或者Shared
cudaDeviceSetCacheConfig(cudaFuncCache cacheConfig) 全局偏好
cudaFuncSetCacheConfig(const void* func, cudaFuncCache cacheConfig) 内核局部偏好
L1 Cache 适用数据访问不规则的任务，Shared Memory适合数据访问规则，线程间通信等任务

RWStructuredBuffer<int> outputBuffer;
groupshared int sharedVar; // 声明一个线程组共享的int

[numthreads(64, 1, 1)]
void CSMain(uint3 globalIdx : SV_DispatchThreadID,
            uint3 localIdx : SV_GroupThreadID,
            uint3 groupIdx : SV_GroupID) {
    if (localIdx.x == 0)
        sharedVar = groupIdx.x; // 线程0写入

    GroupMemoryBarrierWithGroupSync(); // 同步组内所有线程

    outputBuffer[globalIdx.x] = sharedVar; // 所有线程读取
}

constant memory

constant memory 位于显存中，但有专用的Constant Cache加速访问。
Constant Cache与L1 Cache同级，但是独立的缓存单元
每个SM的Constant Cache 8KB，逻辑上总的Constant Memory是64KB
只读
Push Constant 和 UBOS（GLSL）/ cbuffer（HLSL）
经验来讲，大部分需要访问的常量，包括MVP矩阵，灯光参数，渲染设置，骨骼矩阵都放在constant buffer里面，因此其访问相当频繁。实际上，该帧中其指令数仅此于几项高频计算。

cbuffer matrixBuffer : register(b0) { // 绑定到寄存器槽 b0
    float4x4 worldMatrix;
    float4x4 viewMatrix;
    float4x4 projectionMatrix;
};

UDP

Uniform Data Path，处理整个Warp都相同的值时，应用的一种优化，通常Sample/Instruction比较高，说明这里单条指令的消耗更高。当然这里还需要更多的观察

ADU

Address Divergence Unit，地址分歧处理单元。

处理分支和跳转：处理if-else、switch、循环等导致的线程束内执行路径分歧（Warp Divergence）。

处理索引常量加载：负责处理那些索引值在warp内各线程间不一致的常量内存访问。

处理屏障与同步：参与处理__syncthreads()等线程块内同步操作的控制逻辑

可能可以看出分支索引导致的低效问题，但大部分情况下不是主要瓶颈。

SCH

Scheduler 调度器。记录的是 GPU 的调度器和控制逻辑单元处于活跃状态的时间占比。

值得注意的是，空转等待也会被记录在SCH的Sample里面，所以其异常可能反映了调度器等待问题。可能来自：

SM上活跃的 warp 数量很少时，当一些 warp 执行 Load 时（LSU/LDCU的工作），SCH本来应该切另一批warp上来工作，但由于warp不足无法切换，因此会插入 NOP 导致空转。
__syncthreads()，使同一个block内部的warp等待到最后一个执行完毕，引入同步开销，这里也算是空转

CBU

Convergence Barrier Unit，关注线程束分歧问题本身。相比之下，ADU更关注的是一些行为造成的，非合并的地址访问模式问题。只是导致两者的这些行为相近，因此这两个指标也存在一定关联。

Branch：分，即代表分支指令（if-else、switch）本身带来的开销

SIMT：合，即代表__syncthreads() 这类同步等待开销
EXIT：终，即代表Kernel Exit这个结尾段的开销

XU

特殊函数与类型转换单元，类型转换一般在Operation列里会直接写明

TEX

首先，在Shader中我们接触到的图像对象，可以被分为Surface和Texture

Texture，即传统意义上的纹理，只读，且只能用于sample（在Nsight graphics里面叫Fetch），因此通常伴随着硬件滤波（如双线性、三线性插值）
Surface，像是ps中写入render target，或者cs中写入UAV，像素之间通常是独立的，且大部分情况下是数据的载体。有Fetch，Store和Atomic三种操作。Atomic一般通过InterlockedAdd, InterlockedCompareExchange这些函数实现，保证UAV资源写入的原子性，不会因为竞争产生错误。

RTCORE

顾名思义，指会调用Ray tracing core的行为，类似trace ray这种函数。

2. 指标阅读

在所有这个栏目的数据中，都含有 Sample 和 Instruction两项，为了真正理解数据的含义，首先需要对这两项指标做一下严谨的定义。

1. Sample

我们假设，nsight graphics每一秒都对GPU进行大量的sample，随机的sample在数量够多时反映概率。因此，sample指标反映了一段时间内，GPU在进行某项行为的概率。在实际使用时，可以被近似的视为，某项行为在一个区间内的时延情况。

在实际测量中，观察到总Sample数并不是时间均匀的。比如UE render中，在TSR（后处理阶段的时序超分）中，其在322微秒内，总采样次数达到了14536，Lumen Screen Probe Gather 和Direct Lighting也很频繁。但在Nanite draw geometry段就要低的多。从观察上看，似乎与SM Warp Occupancy相关。猜测是，虽然单位时间的sample数一定，但只有sample到活跃核心的才会被保留到报告中。因此workload比较高，sample才比较密集。当然这里目前还没有数值化的证据支撑。

2. Instruction

如果说Sample反映了时延，那么Instruction就对应了频率，即在单位时间内这个行为发生了多少次。配合时延，假如一个行为的Sample很高，但Instruction很低，说明其可能是单次消耗很高，或者流程中存在大量stall的情况，是需要额外警惕的情况。

3. Sample 细节

在Sample中，我们可以看待其被划分为以下这些细节统计项。这些数据是 Warp Stall Reasons。在某一个时刻，一个warp总是在等待什么：可能是在等数据，等执行，等指令，或者等同步。因此，把这些数据全部列出，有助于明确某一项具体的指令的瓶颈在哪里。

Data Dependency

首先是最常见的stall reason，即等数据，其被分为下面两种：

Long Scoreboard (长记分板)：线程束正在等待来自片外内存（如全局内存、纹理内存）的长延迟数据

Short Scoreboard (短记分板)：线程束正在等待来自片上存储（如共享内存）的短延迟数据

可能会有这样的疑问：对于FMA的FP32来说，其本身代表着对两个寄存器的操作，那么为什么其Sample中，Long/Short Scoreboard仍然占据了大部分？理论上不是应该在LSU/LDCU这种显存操作中才比较明显吗？

实际上，这是因为 FMA 依赖的寄存器，在launch这一条指令的时候，不一定准备好了。举一个例子：

float val = tex2D(myTex, uv).x;  // 全局纹理访问，延迟~200-300 cycles
float result = val * 2.0 + 0.5;  // FMA依赖val

阶段一：发射 TEX 指令

调度器选中该 Warp，检查 TEX 指令的源寄存器（纹理坐标）已就绪。

调度器将 TEX 指令发射给 TEX 单元（TMU）。此时，这条指令就结束了，它不再占用调度器的发射带宽。

TEX 单元发送内存请求到 L1/L2/DRAM。

阶段二：Warp 进入“记分板等待”状态（Long Scoreboard）

现在，指令指针（PC）指向下一条指令，即依赖于 TEX 结果的 FMA。

调度器会检查这条 FMA 指令的源寄存器（用于存放纹理返回值的目标寄存器）。记分板显示该寄存器处于 “未就绪（Pending）” 状态。

关键点：由于源操作数未就绪，调度器不会发射这条 FMA。它不会将 FMA 推进 FP 流水线。它只是让这个 Warp 进入 Stall（停顿） 状态，并记录下 Long Scoreboard 停顿原因。

PS：此时实际上TEX单元自己也在记录 Long Scoreboard，也就是说，TEX在等访存，FMA在等TEX写入寄存器，实际上这两个 Long Scoreboard 在绝对时间轴上的重合的。

阶段三：延迟隐藏

根据GPU的设计，此时调度器会调另外一个 Ready 的 Warp 上来，发射它的算术/访存指令

阶段四：数据返回 & 发射 FMA

经过几百个周期后，纹理数据从显存返回，写入目标寄存器。记分板将该寄存器标记为 Ready。

在下一个调度周期，调度器再次检查这个 Warp 的 FMA 指令，发现源寄存器就绪，于是现在才真正将 FMA 指令发射给 FP 流水线执行。

Pipeline Busy

这类停顿源于执行指令所需的硬件单元正忙。

Math Pipe Throttle (数学流水线节流)：指令所需的计算流水线（如ALU, FMA, Tensor Core）正忙。这通常意味着计算任务本身过重。这个指标只在计算相关的 Pipe 中出现
MIO Throttle (MIO 节流)：指令所需的MIO（Memory Input/Output）单元队列已满。MIO负责处理共享内存、特殊数学指令和动态分支等操作。这个指标只在LSU，LDCU这些 Pipe 中出现

Scheduling & Fetch

这类停顿与指令的调度和获取有关。

Dispatch Stall (调度停顿)：因调度资源不可用而无法发射指令。可能是由于指令缓存未命中、流水线冲刷等原因。

No Instructions (无指令)：线程束没有要执行的指令。可能因为指令缓存未命中，或内核本身工作量过小。

Not Selected (未被选中)：线程束本身可以执行（数据已就绪），但调度器选择了其他线程束。

Branch Resolving (分支解析)：线程束正在处理一个分支指令，导致停顿。通常与动态分支或线程束内分歧（Divergence）有关。

Others

Wait：一般指同步等待状态，如__syncthreads()
Selected：正常的工作状态，不是停顿原因
Drain：指流水线或队列正在被“排空”。例如，在调度中间发生的停顿，可能与子通道切换有关