NVIDIA Nsight Compute 入门：从 Copy/Add Kernel 看懂 CUDA Kernel 性能瓶颈

NVIDIA Nsight Compute 入门：从 Copy/Add Kernel 看懂 CUDA Kernel 性能瓶颈

最近在系统学习 CUDA 底层原理和性能优化，正好使用 NVIDIA Nsight Compute 对一个简单 CUDA 程序做了性能分析。

很多刚开始学习 CUDA Profiling 的同学容易混淆两个工具：

Nsight Systems：看整个程序的时间线
Nsight Compute：看单个 CUDA Kernel 内部的性能细节

也就是说：

Nsight Systems 更像“全局时间线分析工具”
Nsight Compute 更像“Kernel 体检报告”

本文通过一个简单的 CUDA Demo，结合 Nsight Compute 的分析结果，系统讲清楚：

1. Nsight Compute 主要看哪些指标
2. Duration、Compute Throughput、Memory Throughput 如何理解
3. Grid Size / Block Size 和源码如何对应
4. Tail Effect 是什么
5. 为什么 Demo Kernel 很短时，GPU 往往吃不满
6. 如何根据 Nsight Compute 的提示做下一步优化

---

一、Nsight Compute 和 Nsight Systems 的区别

在 CUDA 性能分析中，NVIDIA 最常用的两个工具是：

Nsight Systems
Nsight Compute

它们的定位不同。

1. Nsight Systems：看整个程序时间线

Nsight Systems 主要用于分析整个应用的执行流程，例如：

CPU 线程
CUDA API 调用
cudaMalloc
cudaMemcpy
Kernel Launch
cudaDeviceSynchronize
cudaFree
GPU 队列
CPU 等待 GPU 的时间

它适合回答这类问题：

程序整体慢在哪里？
是 cudaMemcpy 慢？
是 cudaMalloc 慢？
是 Kernel 慢？
还是 CPU 在等待 GPU？

简单理解：

Nsight Systems = 程序级时间线分析工具

---

2. Nsight Compute：看单个 Kernel 内部性能

Nsight Compute 专门分析某一个 CUDA Kernel 内部的性能，例如：

SM 利用率
Memory Throughput
Compute Throughput
Occupancy
Warp Stall
L1 / L2 Cache
Register 使用量
Shared Memory 使用量
Instruction 执行效率
Memory Access Pattern

它适合回答这类问题：

这个 Kernel 为什么慢？
是算力没打满？
是显存带宽没打满？
是访存不连续？
是 Occupancy 太低？
是寄存器太多？
是 Warp Stall 太严重？

简单理解：

Nsight Compute = Kernel 级性能体检报告

---

二、Nsight Compute 中常见指标说明

在 Nsight Compute 的 Kernel 列表中，常见指标包括：

Estimated Speedup
Function Name
Demangled Name
Duration
Runtime Improvement
Compute Throughput
Memory Throughput
Registers
Threads
Grid Size
Block Size

下面逐个解释。

---

1. Estimated Speedup

Estimated Speedup = Nsight Compute 估算的理论优化空间

它不是说程序已经提升了多少，而是说：

如果解决当前 Nsight Compute 识别出的主要瓶颈，
理论上可能还能提升多少性能。

例如：

Estimated Speedup = 33.3%

表示 Nsight Compute 认为当前 Kernel 可能存在一个性能问题，如果解决，理论上最多可能提升约 33.3%。

注意：

这是估算值，不是绝对值。

它主要用于帮助我们判断：

当前最值得关注的性能问题是什么。

---

2. Function Name / Demangled Name

这两个字段对应 CUDA Kernel 函数名。

例如源码中有：

Copy<<<grid, block>>>(d_input, d_output);

Add<<<grid, block>>>(d_input, d_output, d_add_output);

那么 Nsight Compute 中通常会看到：

Copy
Add

如果是 C++ 模板函数，编译后函数名可能会被 C++ name mangling 处理，Demangled Name 会显示更容易阅读的原始函数名。

---

3. Duration

Duration = Kernel 本身在 GPU 上执行的时间

例如：

Copy Duration = 3.5 us
Add  Duration = 4.58 us

单位是微秒：

1 ms = 1000 us

所以：

3.5 us  = 0.0035 ms
4.58 us = 0.00458 ms

如果 Kernel 只有几微秒，说明这个 Kernel 本身非常短。

这时候需要注意：

程序整体慢，不一定是 Kernel 慢。

很可能时间花在：

CUDA 初始化
cudaMalloc
cudaMemcpy
cudaDeviceSynchronize
cudaFree
CPU 打印数据
Profiling 工具开销

---

4. Compute Throughput

Compute Throughput = GPU 计算单元利用率

它反映 Kernel 对 GPU 计算资源的使用情况。

如果 Compute Throughput 很低，通常说明：

CUDA Core / Tensor Core 没有被充分利用

可能原因包括：

1. Kernel 计算量太小
2. 每个线程做的计算太少
3. Kernel 被访存拖慢
4. Block 数量太少，SM 吃不满
5. Warp Stall 比较严重

---

5. Memory Throughput

Memory Throughput = 内存通路利用率

它反映 Kernel 对显存、L2 Cache、内存访问通路的使用情况。

如果一个 Kernel：

Compute Throughput 低
Memory Throughput 相对较高

通常说明它更偏向：

Memory-bound

也就是：

不是 GPU 算不过来，而是主要在等数据读写。

典型例子：

out[i] = a[i] + b[i];

这种 Kernel 每个线程只做一次加法，但是需要读写 global memory，所以它很容易成为 memory-bound Kernel。

---

6. Registers / Threads

Registers per Thread = 每个线程使用的寄存器数量

寄存器越多，每个 SM 能同时驻留的 Block / Warp 可能越少，从而影响 Occupancy。

例如：

Registers per Thread = 16

这个数值不算高。

如果一个 Kernel 每个线程使用几十甚至上百个寄存器，就需要关注是否因为寄存器压力过大导致 Occupancy 下降。

---

7. Grid Size / Block Size

这两个指标直接对应 CUDA Kernel Launch 配置：

Kernel<<<grid, block>>>();

例如：

dim3 grid(1024);
dim3 block(256);

Copy<<<grid, block>>>(...);

对应 Nsight Compute 中：

Grid Size  = 1024
Block Size = 256

也就是：

一共启动 1024 个 Block
每个 Block 有 256 个 Thread

总线程数为：

1024 * 256 = 262144 threads

---

三、实验代码说明

下面是本文分析使用的 CUDA Demo。

这个程序主要做了几件事：

1. 在 CPU 侧生成随机输入数据
2. 使用 cudaMalloc 分配 GPU 显存
3. 使用 cudaMemcpy 将数据从 Host 拷贝到 Device
4. 启动 Copy Kernel
5. 启动 Add Kernel
6. 使用 cudaMemcpy 将结果从 Device 拷贝回 Host
7. 使用 cudaEvent 统计整体耗时
8. 使用 NVTX 标记不同阶段，方便 Nsight 工具观察

核心代码如下：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <thread>
#include <chrono>
#include <stdio.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <iostream>
#include <vector>
#include <random>
#include <chrono>
#include <cmath>

#include "utils.h"
#include "kernel_function.h"

#include <nvtx3/nvToolsExt.h>

int run_cuda()
{
    const int blockSize = 256;

    // RTX 5080 示例：SM 数量约为 84
    // 每个 SM 最多可驻留多个 block
    const int numBlocks = 1024;

    const int N = numBlocks * blockSize;

    printf("blockSize=%d, numBlocks=%d, N=%d\n", blockSize, numBlocks, N);

    // ---------- 1. 生成随机输入数据 ----------
    std::vector<float> h_input(N);

    std::random_device rd;
    std::mt19937 gen(rd());
    std::uniform_real_distribution<float> dist(-100.0f, 100.0f);

    for (int i = 0; i < N; ++i)
    {
        h_input[i] = dist(gen);
    }

    ShowArray(h_input, "input_data");

    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    cudaEventRecord(start);

    // ---------- 2. 分配设备内存 ----------
    float* d_input;
    float* d_output;
    float* d_add_output;

    nvtxRangePushA("cudaMalloc");
    CUDA_CHECK(cudaMalloc(&d_input, N * sizeof(float)));
    CUDA_CHECK(cudaMalloc(&d_output, numBlocks * sizeof(float)));
    CUDA_CHECK(cudaMalloc(&d_add_output, N * sizeof(float)));
    nvtxRangePop();

    // ---------- 3. Host to Device ----------
    nvtxRangePushA("H2D cudaMemcpy");
    CUDA_CHECK(cudaMemcpy(d_input, h_input.data(), N * sizeof(float), cudaMemcpyHostToDevice));
    nvtxRangePop();

    // ---------- 4. Kernel Launch ----------
    dim3 grid(numBlocks);
    dim3 block(blockSize);

    nvtxRangePushA("Copy Kernel");
    Copy<<<grid, block>>>(d_input, d_output);
    cudaDeviceSynchronize();
    nvtxRangePop();

    nvtxRangePushA("Add Kernel");
    Add<<<grid, block>>>(d_input, d_output, d_add_output);
    cudaDeviceSynchronize();
    nvtxRangePop();

    // ---------- 5. Device to Host ----------
    std::vector<float> h_output(numBlocks);
    std::vector<float> h_add_output(N);

    nvtxRangePushA("D2H cudaMemcpy");

    CUDA_CHECK(cudaMemcpy(
        h_output.data(),
        d_output,
        numBlocks * sizeof(float),
        cudaMemcpyDeviceToHost
    ));

    ShowArray(h_output, "cuda_output_data");

    CUDA_CHECK(cudaMemcpy(
        h_add_output.data(),
        d_add_output,
        N * sizeof(float),
        cudaMemcpyDeviceToHost
    ));

    nvtxRangePop();

    ShowArray(h_add_output, "Add_cuda_output_data");

    // ---------- 6. 释放资源 ----------
    cudaFree(d_input);
    cudaFree(d_output);
    cudaFree(d_add_output);

    cudaEventRecord(stop);
    cudaEventSynchronize(stop);

    float ms;
    cudaEventElapsedTime(&ms, start, stop);

    std::cout << "Average kernel time: " << ms << " ms\n";
    printf("=======================================================\n");

    return 0;
}

int main(int argc, char* argv[])
{
    run_cuda();
}

---

四、Nsight Compute 抓到的是哪部分？

需要特别注意：

Nsight Compute 默认主要分析 CUDA Kernel 本身。

也就是说，它重点看的是：

Copy<<<grid, block>>>(d_input, d_output);

Add<<<grid, block>>>(d_input, d_output, d_add_output);

而不是重点看：

cudaMalloc
cudaMemcpy
cudaFree
cudaDeviceSynchronize

这些 API 更适合用 Nsight Systems 分析。

所以，当 Nsight Compute 界面中看到两行 Kernel：

0    Copy
1    Add

它们对应的正是源码中的：

Copy<<<1024, 256>>>(...);
Add <<<1024, 256>>>(...);

---

五、Kernel 配置和源码的对应关系

源码中配置如下：

const int blockSize = 256;
const int numBlocks = 1024;

dim3 grid(numBlocks);
dim3 block(blockSize);

因此实际 Kernel Launch 是：

Copy<<<1024, 256>>>(d_input, d_output);

Add<<<1024, 256>>>(d_input, d_output, d_add_output);

对应 Nsight Compute 中：

Grid Size  = 1024
Block Size = 256

总线程数量为：

1024 * 256 = 262144

也就是：

N = 262144 个 float

数据量大约是：

262144 * 4 Byte = 1 MB

对于 RTX 5080 这种级别的 GPU 来说，1 MB 数据量非常小。

所以这个 Demo 的核心问题不是“GPU 算力不够”，而是：

Kernel 太短，任务规模太小，GPU 没有被充分喂满。

---

六、两个 Kernel 的耗时分析

Nsight Compute 中可以看到两个 Kernel 的执行耗时：

Kernel	Duration	换算
Copy	3.5 us	0.0035 ms
Add	4.58 us	0.00458 ms

这说明：

Copy Kernel 和 Add Kernel 都非常短。

如果程序整体运行时间看起来比较长，不要马上判断是 Kernel 慢。

因为整体耗时可能包括：

1. CUDA Runtime 初始化
2. cudaMalloc 显存分配
3. cudaMemcpy 数据拷贝
4. cudaDeviceSynchronize 同步等待
5. cudaFree 释放显存
6. CPU 侧 ShowArray 打印
7. Nsight Profiling 工具本身的采样开销

所以优化 CUDA 程序时，一定要区分：

Kernel 本身耗时
CUDA API 耗时
CPU 等待耗时
Host / Device 数据拷贝耗时
Profiling 开销

这也是为什么实际工程中通常需要：

Nsight Systems + Nsight Compute 结合使用

---

七、Compute Throughput 和 Memory Throughput 分析

Nsight Compute 中两个 Kernel 的关键指标大致如下：

Kernel	Compute Throughput	Memory Throughput	初步判断
Copy	20.5%	31.54%	偏访存，计算量很小
Add	13.08%	48.83%	更明显的 memory-bound

---

1. Copy Kernel

Copy Kernel 的特点是：

读 global memory
写 global memory
几乎没有复杂计算

因此它本质上就是一个访存型 Kernel。

如果 Compute Throughput 不高，这是正常的，因为它根本没有太多计算任务。

---

2. Add Kernel

Add Kernel 的特点一般是：

out[i] = input[i] + something;

或者类似的逐元素加法。

这种 Kernel 每个线程只做很少的数学计算，但是需要访问 global memory。

所以它的典型特征是：

Compute Throughput 低
Memory Throughput 相对更高

这说明它更偏向：

Memory-bound Kernel

也就是说：

GPU 不是算不过来，而是主要在等数据读写。

对于这种 Kernel，单纯优化数学计算通常意义不大。

更应该关注：

1. 访存是否连续
2. 是否合并访问 coalesced access
3. 是否存在不必要的 global memory 读写
4. 是否可以减少中间结果写回
5. 是否可以合并多个小 Kernel
6. 是否可以提升数据规模，让 GPU 吃满

---

八、为什么 Estimated Speedup 是 33.3%？

Nsight Compute 中可能会提示：

Estimated Speedup = 33.3%

并且主要原因是：

Tail Effect

这里的 33.3% 不是说程序已经优化了 33.3%，而是说：

Nsight Compute 认为当前 Kernel 存在 Tail Effect，
如果这个问题被缓解，理论上可能有一定性能提升空间。

---

九、Tail Effect 是什么？

Tail Effect 可以理解为：

GPU 执行到最后一批 block 时，剩余 block 数太少，
导致大量 SM 空闲，GPU 利用率下降。

假设：

RTX 5080 有 84 个 SM
每个 SM 可同时驻留约 6 个 block

那么一个完整 wave 大约可以容纳：

84 * 6 = 504 个 block

当前程序启动了：

numBlocks = 1024

因此执行过程大致是：

第 1 wave：504 blocks
第 2 wave：504 blocks
第 3 wave：16 blocks

也就是：

504 + 504 + 16 = 1024

前两波 GPU 基本是满的。

但最后一波只剩下 16 个 block。

这意味着：

只有少量 SM 有活干，大量 SM 处于空闲状态。

这就是 Tail Effect。

---

Tail Effect 图解

完整 wave：

SM0    block block block block block block
SM1    block block block block block block
SM2    block block block block block block
...
SM83   block block block block block block

最后 partial wave：

SM0    block
SM1    block
SM2    block
...
SM15   block
SM16   idle
SM17   idle
...
SM83   idle

最后一批任务太少，GPU 吃不满。

所以 Nsight Compute 会提示：

Tail Effect

---

十、Tail Effect 对这个 Demo 严重吗？

对这个 Demo 来说，Tail Effect 不是严重问题。

原因是：

当前数据量太小
Kernel 执行时间太短
总数据只有约 1 MB

源码中：

const int blockSize = 256;
const int numBlocks = 1024;
const int N = numBlocks * blockSize;

所以：

N = 1024 * 256 = 262144

float 类型是 4 字节：

262144 * 4 = 1048576 Byte ≈ 1 MB

对 RTX 5080 这种 GPU 来说，1 MB 数据非常小。

因此：

Kernel Launch 开销
cudaDeviceSynchronize 开销
Profiling 开销
Tail Effect

都会被放大。

这个阶段不要急着说 Kernel 写得不好，而应该先判断：

当前 Demo 太小，GPU 没吃饱。

---

十一、如何减少 Tail Effect？

方法一：增加 numBlocks

当前配置：

const int numBlocks = 1024;

可以改成：

const int numBlocks = 10080;

因为：

10080 = 504 * 20

这意味着大约可以形成 20 个完整 wave。

修改后：

const int blockSize = 256;
const int numBlocks = 10080;
const int N = numBlocks * blockSize;

此时：

N = 10080 * 256 = 2580480

数据量大约为：

2580480 * 4 Byte ≈ 10 MB

相比原来的 1 MB，更适合做 Profiling。

---

方法二：让 Block 数接近 wave size 的整数倍

当前：

1024 = 504 * 2 + 16

最后剩余 16 个 block，尾部效应明显。

可以尝试：

1008  = 504 * 2
1512  = 504 * 3
5040  = 504 * 10
10080 = 504 * 20

例如：

const int numBlocks = 1008;

这样最后不会留下很小的 partial wave。

不过需要注意：

真实业务中不一定为了消除 Tail Effect 而强行调整数据规模。

这个方法更适合 Benchmark 和学习实验。

---

方法三：加大整体数据量

如果数据量足够大，即使最后有 partial wave，它占总时间的比例也会变小。

例如：

const int numBlocks = 1024 * 100;

这样即使最后仍然存在尾部不满，影响也会被摊薄。

---

十二、L2 Slices Workload Imbalance 是什么？

Nsight Compute 还可能提示：

L2 Slices Workload Imbalance

意思是：

不同 L2 Cache slice 的工作量不均衡。

简单理解：

线程访问 global memory
        ↓
请求进入 L2 Cache
        ↓
L2 Cache 分成多个 slice
        ↓
有的 slice 压力大，有的 slice 比较闲

这通常和 global memory 访问分布有关。

不过在这个 Demo 中，Kernel 太小，数据量太少，所以这个提示暂时不是第一优先级。

当前优化优先级应该是：

第一优先级：确认 Kernel 是否太小
第二优先级：理解 Tail Effect
第三优先级：观察访存模式是否合理
第四优先级：再深入分析 L2 imbalance

---

十三、Registers per Thread = 16 如何理解？

Nsight Compute 中显示：

Registers per Thread = 16

这表示：

每个线程使用 16 个寄存器

这个数值不高。

一般来说：

寄存器使用过多
        ↓
每个 SM 能同时驻留的 warp / block 变少
        ↓
Occupancy 下降
        ↓
隐藏访存延迟能力下降

但是当前 Kernel 只有 16 个 registers/thread，所以：

寄存器压力不是主要瓶颈。

也就是说，这个 Demo 的主要问题不是：

寄存器太多

而是：

Kernel 太小
任务量太少
尾部效应明显
访存型 Kernel 计算量不足

---

十四、当前两个 Kernel 的性能结论

可以总结如下：

Copy Kernel：
    Duration ≈ 3.5 us
    Compute Throughput ≈ 20.5%
    Memory Throughput ≈ 31.54%
    特点：访存型 Kernel，计算量很少
    问题：Kernel 太短，Tail Effect 明显

Add Kernel：
    Duration ≈ 4.58 us
    Compute Throughput ≈ 13.08%
    Memory Throughput ≈ 48.83%
    特点：更明显的 memory-bound
    问题：访存占比高，Kernel 太小，GPU 没吃满

核心判断：

当前不是 GPU 算力不够，
而是任务规模太小，Kernel 太短，GPU 没有被充分喂满。

---

十五、源码和 Nsight Compute 的对应关系

可以用下面这张逻辑图理解：

CPU run_cuda()
    │
    ├── cudaMalloc
    │       Nsight Compute 默认不重点分析
    │       更适合用 Nsight Systems 看
    │
    ├── cudaMemcpy H2D
    │       Host → Device 数据拷贝
    │       Nsight Compute 默认不重点分析
    │
    ├── Copy<<<1024, 256>>>
    │       Nsight Compute 第一行：Copy
    │       Duration: 约 3.5 us
    │       Grid Size: 1024
    │       Block Size: 256
    │       Registers/thread: 16
    │       主要问题：Kernel 太小 + Tail Effect
    │
    ├── cudaDeviceSynchronize
    │       CPU 等待 GPU 执行完成
    │
    ├── Add<<<1024, 256>>>
    │       Nsight Compute 第二行：Add
    │       Duration: 约 4.58 us
    │       Grid Size: 1024
    │       Block Size: 256
    │       Registers/thread: 16
    │       主要问题：Memory-bound + Kernel 太小 + Tail Effect
    │
    ├── cudaDeviceSynchronize
    │       CPU 再次等待 GPU
    │
    ├── cudaMemcpy D2H
    │       Device → Host 数据拷贝
    │       Nsight Compute 默认不重点分析
    │
    └── cudaFree
            显存释放
            更适合用 Nsight Systems 看

---

十六、下一步实验建议

为了让 Nsight Compute 的结果更有参考价值，可以先做下面几个实验。

---

1. 增大 numBlocks

将：

const int numBlocks = 1024;

改成：

const int numBlocks = 10080;

完整配置：

const int blockSize = 256;
const int numBlocks = 10080;
const int N = numBlocks * blockSize;

这样数据量从约 1 MB 增加到约 10 MB。

预期现象：

Kernel Duration 变长
Tail Effect 影响下降
Compute / Memory Throughput 更稳定
Nsight Compute 报告更有参考意义

---

2. 注释掉 ShowArray

当前代码中有：

ShowArray(h_input, "input_data");
ShowArray(h_output, "cuda_output_data");
ShowArray(h_add_output, "Add_cuda_output_data");

这些 CPU 打印会严重影响整体程序耗时。

做性能分析时建议先注释掉：

// ShowArray(h_input, "input_data");
// ShowArray(h_output, "cuda_output_data");
// ShowArray(h_add_output, "Add_cuda_output_data");

否则很容易误判：

程序慢是因为 CUDA 慢

但实际上可能是：

CPU 打印太慢

---

3. 减少不必要的 cudaDeviceSynchronize

当前每个 Kernel 后面都调用了：

cudaDeviceSynchronize();

这会强制 CPU 等待 GPU 完成。

在学习阶段这样写有助于调试，但是在真实工程中，频繁同步会破坏异步执行能力。

如果不是必须立即拿到结果，可以考虑：

减少同步点
使用 cudaEvent 计时
使用 Stream 构建异步流水线

---

4. 使用 Nsight Systems 看整体耗时

Nsight Compute 主要看 Kernel 内部。

如果想知道：

cudaMalloc 是否慢
cudaMemcpy 是否慢
CPU 是否在等待 GPU
Kernel Launch 间隔是否过大

应该使用：

Nsight Systems

建议组合方式：

第一步：用 Nsight Systems 看程序整体时间线
第二步：找到最耗时的 Kernel
第三步：用 Nsight Compute 深入分析该 Kernel 内部瓶颈

这才是比较完整的 CUDA Profiling 流程。

---

十七、实际工程中的优化方向

对于这个 Demo，可以考虑的优化方向包括：

1. 增大数据规模，让 GPU 吃满
2. 减少小 Kernel 数量，尝试 Kernel Fusion
3. 减少不必要的 cudaDeviceSynchronize
4. 将 cudaMalloc / cudaFree 移出循环
5. 使用 Memory Pool 复用显存
6. 使用 pinned memory 提升 H2D / D2H 拷贝效率
7. 使用 cudaMemcpyAsync + Stream 做异步流水线
8. 检查 global memory 是否连续访问
9. 避免不必要的 global memory 中间写回
10. 用 Nsight Systems + Nsight Compute 联合分析

在 AI 推理框架中，例如 TensorRT、vLLM、TensorRT-LLM，通常不会频繁调用：

cudaMalloc();
cudaFree();

而是更倾向于：

预分配显存
Memory Pool
KV Cache 复用
CUDA Graph
Stream 并发
Kernel Fusion
算子融合

原因就是：

频繁的显存申请、释放、同步和小 Kernel Launch，
都会让 GPU 无法持续满负载运行。

---

十八、面试中如何回答 Nsight Compute？

如果面试官问：

你怎么用 Nsight Compute 分析 CUDA Kernel？

可以这样回答：

我一般先用 Nsight Systems 看整个程序时间线，判断瓶颈是在 CPU、cudaMemcpy、cudaMalloc、同步等待，还是 Kernel 本身。

如果确认是 Kernel 慢，再用 Nsight Compute 进入 Kernel 内部分析。

在 Nsight Compute 中，我会重点看 Duration、Compute Throughput、Memory Throughput、Occupancy、Registers per Thread、Shared Memory、Warp Stall、L2 Cache、Memory Access Pattern 等指标。

如果 Compute Throughput 低而 Memory Throughput 高，通常说明 Kernel 偏 memory-bound；如果两者都低，可能是 Kernel 太小、Occupancy 不够、Tail Effect 或同步开销明显。

以 Copy/Add 这种简单 Kernel 为例，Duration 只有几微秒，Compute Throughput 不高，Memory Throughput 相对更明显，说明它主要不是计算瓶颈，而是任务规模太小、访存占比高、GPU 没有吃满。

这个回答比简单说“调 Block Size”更专业。

因为 CUDA 优化的正确流程是：

先 Profile
再定位瓶颈
最后针对性优化

而不是一上来就盲目改参数。

---

十九、本文核心总结

这次 Nsight Compute 分析可以总结为：

1. Nsight Compute 主要分析单个 CUDA Kernel 内部性能
2. Nsight Systems 更适合看整个程序时间线
3. Copy 和 Add 对应源码中的两个 Kernel Launch
4. Grid Size = 1024，Block Size = 256，与 <<<grid, block>>> 完全对应
5. Copy Kernel 约 3.5 us，Add Kernel 约 4.58 us，Kernel 本身非常短
6. Compute Throughput 低，说明计算单元没有打满
7. Add 的 Memory Throughput 更高，说明它更偏 memory-bound
8. Estimated Speedup 33.3% 主要来自 Tail Effect
9. Tail Effect 是最后一批 block 太少，导致部分 SM 空闲
10. 当前 Demo 的主要问题不是代码错误，而是任务规模太小，GPU 没吃饱
11. 做 Profiling 时应增大数据量，并注释掉 CPU 打印
12. 实际工程中应结合 Nsight Systems 和 Nsight Compute 一起分析

最后记住一句话：

Nsight Compute 不是告诉你“代码哪里错了”，
而是告诉你“Kernel 的硬件资源到底有没有被充分利用”。

CUDA 优化的核心不是猜，而是：

Profile → 定位瓶颈 → 针对性优化 → 再次验证