参考：
https://zhuanlan.zhihu.com/p/707107808
BBuf
官方文档

Nsight Compute简介

Nsight Compute是一个CUDA kernel分析器，它通过硬件计数器和软件收集指标。它使用内置的专业知识来检测kernel常见的性能问题并指出发生这些问题的位置并给出一些解决方法的建议。这一内置规则集和指南就是我们所说的Guided Analysis。下面就结合Lecture1的例子来深入了解下Nsight Compute。

在Nsight Compute中，如果我们把鼠标悬停在各个指标上，我们能获得对应的讲解。

Nsight Compute Profile流程

BBuf的Blog中使用的是 Triton 实现的矩阵开方代码使用Nsight Compute进行Profile，我们这里使用naive方式的sgemm和cublas的gemm进行实验。打开方式是直接使用管理员权限的Nsight Compute，完全不使用用ncu 命令行。实验环境为Windows专业版，显卡为GTX 1660Ti。
下面给出naive的kernel代码。

__global__ void naiveSgemm(float *__restrict__ a, float *__restrict__ b,  
                           float *__restrict__ c, const int M, const int N,  
                           const int K) {  
  
  int n = blockIdx.x * blockDim.x + threadIdx.x;  
  int m = blockIdx.y * blockDim.y + threadIdx.y;  
  if (m < M && n < N) {  
    float psum = 0.0;  
#pragma unroll  
    for (int k = 0; k < K; k++) {  
      psum += a[OFFSET(m, k, K)] * b[OFFSET(k, n, N)];  
    }    c[OFFSET(m, n, N)] = psum;  
  }}

cublas的代码：

cublasSgemm(cublas_handle, CUBLAS_OP_N, CUBLAS_OP_N, N, M, K, &cublas_alpha,  
            d_b, N, d_a, K, &cublas_beta, d_c, N);

下面开始分析，注意需要管理员权限打开Nsight Compute

注意Sections选择full

然后直接launch。弹出的terminal上可以看到我们程序的输出。

Summary

Summary部分长这样，可以看到我们启动的三个kernel：第一个是warm up，第二个是cublas的实现，第三个是我们的naive cuda实现。

使用page可以切换不同的界面，我们先看Summary

这部分内容是选中kernel的一些信息，先混个脸熟。

下面是表示这个Summary选择下可以看到的Profile程序里面有哪些Kernel。下面是指标的一些含义，同样现在不需要了解每个的具体含义。

Compute Throughput

在 CUDA 中，Compute Throughput 百分比 是衡量 GPU 计算单元（SM, Streaming Multiprocessor）的利用率的指标。
鼠标悬停在这一指标时，会出现如下提示：

同 SM Throughput，SM（流式多处理器）是 GPU 的基本计算单元，负责并行执行 CUDA 程序中的大量线程。每个 SM 可以并行地执行多个warp，并且每个 warp 包含 32 个线程。每个 SM 有一定数量的计算资源（如 ALU、寄存器、共享内存等），这些资源用于执行程序中的指令。

• Warp 是 32 个线程的组合，它们在同一周期内执行相同的指令。
• CTA（Cooperative Thread Array）或线程块，是由多个 warp 组成的
• SMSPs：每个SM被划分为四个处理块，称为SM子分区。 SM子分区是SM上的主要处理单元。 一个子分区管理一个固定大小的warp池。

计算步骤：
计算 Compute Throughput 百分比涉及以下几个步骤：

Step 1: 获取 GPU 的理论最大 FLOPS
首先，你需要知道 GPU 的规格，包括以下信息：

• SM 核心数量：例如 NVIDIA A100 有 108 个 SM，每个 SM 包含 64 个 CUDA 核心，总共 6912 个 CUDA 核心。
• 时钟频率：例如 A100 的基础时钟频率是 1410 MHz（1.41 GHz）。
• 每周期浮点操作数：每个 CUDA 核心通常可以在每个时钟周期执行 2 次浮点操作（一次 FMA）。

理论 FLOPS 计算公式：

$$ \text{理论 FLOPS} = \text{核心数量} \times 2 \times \text{时钟频率 (Hz)} $$

示例：

• GPU：NVIDIA A100
• 核心数：6912
• 时钟频率：1.41 GHz
• 理论 FLOPS：

$$ 6912×2×1.41×10^9=19.47 TFLOPS $$

Step 2: 计算实际执行的计算操作数（FLOP）
实际的 FLOP 需要通过程序运行时收集，常用的方法包括：

• 使用 CUDA Profiler 工具：
  • Nsight Compute 和 nvprof 可以记录 CUDA 核函数执行的 FLOP 数。
  • FLOP 的具体计数通常包含单精度 (FP32)、双精度 (FP64) 和混合精度 (Tensor Core) 的统计数据。
• 手动估算： 如果你清楚代码中执行的计算逻辑，可以估算每次循环或每个线程的浮点操作数，再乘以线程总数和循环次数。

Step 3: 计算 Compute Throughput%

Memory Throughput

Memory Throughput 百分比 是衡量 CUDA 程序运行时 GPU 内存带宽利用效率的一个关键指标。通过计算 Memory Throughput 百分比，可以评估程序对 GPU 全局内存、共享内存或其他存储资源的访问效率，从而识别内存相关的性能瓶颈并进行优化。

计算内存管道吞吐量 （此吞吐量指标表示在所有子单元实例的经过周期内达到的峰值持续率的百分比）

Registers

寄存器：每个子分区有一组32位寄存器，由硬件以固定大小的块分配。

线程：在GPU的一个SM单元上运行的单个线程。

Cycles

CUDA程序在GPU上运行时，每个指令都需要一定的硬件时钟周期来完成。一个cycle代表一个时钟信号的周期，是GPU硬件执行工作的基本时间单位。
Warp执行与Cycles： GPU的基本调度单位是warp（一组32个线程）。在CUDA中，warp中的所有线程是同步执行的（SIMD模型）。一个指令执行完成需要一定数量的cycles。

延迟与吞吐量： 一些CUDA指令（如内存读取、计算指令）可能需要多个周期（cycles）才能完成。这取决于指令类型、寄存器访问、线程分支情况以及内存访问模式等。

Details

GPU Speed Of Light Throughput

首先是 GPU Speed Of Light Throughput部分，它通常位于Details部分的顶部。它清晰的描述了GPU资源的利用情况。在下面的截图中，我们同样可以通过鼠标悬停的方式去看每个指标的细节。

Balanced Throughput： 平衡吞吐量的提示，此处计算吞吐量与内存吞吐量都约为 77.89%，表明任务优化良好，避免了瓶颈。
Roofline Analysis：当前设备的 fp32（单精度浮点）与 fp64（双精度浮点）性能比为 32:1。此内核只达到了设备 fp32 峰值性能的 10%，说明仍有优化空间。双精度性能（fp64）未被利用（达成 0% 峰值性能）。

下面这些了解一下就行，暂时还没有用到。

Compute Throughput Breakdown

接下来是Roofline

关于roofline的可以看另一个blog，有详细的讲解。
可以看到naiveSgemm的内存访问和计算都有优化的空间。

Compute Workload Analysis

总体指标

分析与优化方向：
内存瓶颈：LSU 利用率达到 78.2%，显著高于其他管道，说明内存加载/存储操作是当前任务的主要性能瓶颈。
优化方向：

• 减少不必要的内存访问，增加数据局部性（使用共享内存或缓存）。
• 合并内存操作，减少内存访存的频率。

计算单元：FMA 和 ALU 管道的利用率尚未达到满负载，说明计算性能尚未完全释放。
优化方向：

• 增加浮点运算和整数运算的并行度。
• 优化核函数（Kernel）代码逻辑，提升计算密集度。

整体利用率：SM Busy 和 Issue Slots Busy 均为 29.42%，表明 GPU 的整体利用率较低。
优化方向：

• 提高任务的线程并行度。
• 减少线程阻塞和数据依赖，提高 GPU 核心的占用率

Memory Workload Analysis

Memory Workload Analysis（内存工作负载分析），用于分析 GPU 的内存资源利用情况，包括内存访问模式、带宽使用情况以及潜在的内存瓶颈。

顶部主要指标

L2 Load Access Pattern: 这部分详细描述了从 L1/TEX 缓存加载到 L2 缓存时的访问模式

当前分析：Sectors per L2 Request 为 1.6，远小于 4。说明内存访问模式未能充分利用 L2 缓存行，每次加载的数据块较小，导致带宽浪费或缓存效率低下。原因可能是 非合并内存访问（Uncoalesced Accesses），即多个线程对不连续的内存地址进行访问。

接下来是Memory Chart：

任务的大部分内存请求都集中在 全局内存，共发起了 268.44M 请求，可能是性能瓶颈的主要来源。
L1 缓存：高命中率（94.93%），有效减少了对全局内存的访问，但仍有部分请求流入 L2 缓存。
L2 缓存：命中率较低（49.22%），需要优化数据访问模式，进一步减少全局内存访问。
内存瓶颈：
全局内存（Device Memory）的流量较高（573.34 MB），可能导致较高的延迟。

Scheduler Statistics

Scheduler Statistics（调度器统计） 的分析结果，用于评估 GPU 的线程调度器在内核执行期间的活跃度和利用率。调度器是 GPU 的核心组件之一，负责管理和调度多个线程组（即 warps）执行指令。

分析与优化

调度器活跃 Warp 数量： Active Warps Per Scheduler 接近理论最大值（7.99 / 8 warp），表明任务在调度器上分配了足够多的 Warp，调度器的硬件资源已接近饱和。

Eligible Warp 数量较少 ： Eligible Warps Per Scheduler 仅为 0.74 warp，说明大部分 Warp 处于停顿状态，可能因为以下原因：
1. 内存延迟：Warp 等待内存访问完成。
2. 数据依赖：Warp 等待前一指令的计算结果。
3. 资源争用：Warp 等待寄存器或其他硬件资源的可用性。

Issued Warp 数量低 ： 每个调度器平均每周期只有 0.29 warp 实际发出了指令，说明大部分调度器的发射槽（Issue Slot）处于空闲状态，未能充分利用 GPU 的并行计算能力。

No Eligible 占比高 ： 调度器中 70.56% 的周期没有可发出指令的 Warp，进一步表明 Warp 停顿严重。

Warp State Statistics

Warp State Statistics（线程束状态统计） 的分析结果，用于了解 GPU 任务执行中 Warp（线程束）的执行和停顿状态。Warp 停顿（Stall）是 GPU 性能的主要瓶颈之一，这个图揭示了 Warp 在每周期中的状态分布以及主要的停顿原因。

柱状图显示了 Warp 状态在所有周期中的分布，包括主要停顿原因和其他状态：

分析与优化

Instruction Statistics

Instruction Statistics（指令统计），主要用来分析 GPU 内核中已执行和已发射的低级汇编指令（SASS）。这些统计信息提供了关于指令类型和频率的见解，帮助了解指令发射和执行的效率以及潜在的性能瓶颈。

Occupancy

Occupancy（占用率） 的相关指标，用于分析 GPU 核函数（Kernel）在 Streaming Multiprocessor（SM）上的 Warp 分配和利用情况。

块大小对性能的影响，每块共享内存用量对性能的影响

Source Counters

在这个电脑上只有SASS的代码，具体可以看这个BLOG CUDA MODE

总结

使用Nsight Compute进行核函数的性能分析，主要可以概括为以下几个步骤：

1. 定位性能瓶颈：通过采集 GPU 执行的详细性能数据，确定性能瓶颈所在（内存、计算、线程分配等）。
2. 分析关键指标：重点分析影响性能的指标，如计算效率、内存利用率、线程并行度、指令吞吐量等。
3. 优化与验证：针对瓶颈采取优化措施，并通过再次分析验证优化效果。

最后，注意Nsight Compute中的提示，通常都是一些可以优化的点。