1. 背景
通过 demo 代码生成 nsys profile 文件，观察 nsys ui 工具的各个部分。
2. 安装 cuda 环境（按需）
# 对于Ubuntu/Debian系统
sudo apt install nvidia-cuda-toolkit
# 对于CentOS/RHEL系统
sudo yum install cuda
# 永久生效可以添加到~/.bashrc
echo 'export PATH=/usr/local/cuda/bin:$PATH' >> ~/.bashrc
source ~/.bashrc
3. 安装 nsys（按需）
ubuntu 环境
（1）查看 cuda 版本
root@9187eedb402d:/vllm-workspace# nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Fri_Feb_21_20:23:50_PST_2025
Cuda compilation tools, release 12.8, V12.8.93
Build cuda_12.8.r12.8/compiler.35583870_0
（2）安装
12.8 对应 cuda 版本。
apt-get install cuda-nsight-systems-12.8
4. demo cuda 代码
（1）使用 CUDA 的 demo 代码
demo cuda 代码
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <curand.h>
#include <curand_kernel.h>
// CUDA核函数：向量加法
__global__ void vectorAdd(const float *a, const float *b, float *c, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}
// CUDA核函数：矩阵乘法
__global__ void matrixMultiply(const float *a, const float *b, float *c, int width) {
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < width && col < width) {
float sum = 0.0f;
for (int k = 0; k < width; k++) {
sum += a[row * width + k] * b[k * width + col];
}
c[row * width + col] = sum;
}
}
// CUDA核函数：初始化随机数
__global__ void initCurandStates(curandState *states, unsigned long long seed) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
curand_init(seed, idx, 0, &states[idx]);
}
// CUDA核函数：生成随机数
__global__ void generateRandomNumbers(curandState *states, float *data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
data[idx] = curand_uniform(&states[idx]);
}
}
int main(int argc, char **argv) {
const int VECTOR_SIZE = 1 << 20; // 1M elements
const int MATRIX_WIDTH = 1024;
const int MATRIX_SIZE = MATRIX_WIDTH * MATRIX_WIDTH;
float *h_a, *h_b, *h_c; // 主机端数据
float *d_a, *d_b, *d_c; // 设备端数据
curandState *d_states; // 随机数状态
int deviceId;
cudaGetDevice(&deviceId);
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, deviceId);
printf("Using device %d: %s\n", deviceId, deviceProp.name);
// 分配主机端内存
h_a = (float*)malloc(VECTOR_SIZE * sizeof(float));
h_b = (float*)malloc(VECTOR_SIZE * sizeof(float));
h_c = (float*)malloc(VECTOR_SIZE * sizeof(float));
// 初始化主机端数据
for (int i = 0; i < VECTOR_SIZE; i++) {
h_a[i] = (float)i;
h_b[i] = (float)(VECTOR_SIZE - i);
}
// 分配设备端内存
cudaMalloc((void**)&d_a, VECTOR_SIZE * sizeof(float));
cudaMalloc((void**)&d_b, VECTOR_SIZE * sizeof(float));
cudaMalloc((void**)&d_c, VECTOR_SIZE * sizeof(float));
// 主机到设备的数据传输
cudaMemcpy(d_a, h_a, VECTOR_SIZE * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, VECTOR_SIZE * sizeof(float), cudaMemcpyHostToDevice);
// 设置向量加法的线程块和网格维度
int threadsPerBlock = 256;
int blocksPerGrid = (VECTOR_SIZE + threadsPerBlock - 1) / threadsPerBlock;
// 启动向量加法核函数
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, VECTOR_SIZE);
// 设备到主机的数据传输
cudaMemcpy(h_c, d_c, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost);
// 验证结果
float sum = 0.0f;
for (int i = 0; i < VECTOR_SIZE; i++) {
sum += h_c[i];
}
printf("Vector add checksum: %.2f\n", sum);
// 矩阵乘法部分
// 重新分配设备端内存用于矩阵
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
cudaMalloc((void**)&d_a, MATRIX_SIZE * sizeof(float));
cudaMalloc((void**)&d_b, MATRIX_SIZE * sizeof(float));
cudaMalloc((void**)&d_c, MATRIX_SIZE * sizeof(float));
// 分配随机数状态内存
cudaMalloc((void**)&d_states, MATRIX_SIZE * sizeof(curandState));
// 初始化随机数生成器
dim3 threads(16, 16);
dim3 blocks(MATRIX_WIDTH / threads.x, MATRIX_WIDTH / threads.y);
initCurandStates<<<blocks, threads>>>(d_states, 1234ULL);
// 生成随机矩阵
generateRandomNumbers<<<blocks, threads>>>(d_states, d_a, MATRIX_SIZE);
generateRandomNumbers<<<blocks, threads>>>(d_states, d_b, MATRIX_SIZE);
// 同步以确保随机数生成完成
cudaDeviceSynchronize();
// 启动矩阵乘法核函数
matrixMultiply<<<blocks, threads>>>(d_a, d_b, d_c, MATRIX_WIDTH);
// 同步以确保矩阵乘法完成
cudaDeviceSynchronize();
// 分配主机端内存用于结果矩阵
float *h_result = (float*)malloc(MATRIX_SIZE * sizeof(float));
// 设备到主机的数据传输
cudaMemcpy(h_result, d_c, MATRIX_SIZE * sizeof(float), cudaMemcpyDeviceToHost);
// 计算结果校验和
float checksum = 0.0f;
for (int i = 0; i < MATRIX_SIZE; i++) {
checksum += h_result[i];
}
printf("Matrix multiply checksum: %.2f\n", checksum);
// 释放内存
free(h_a);
free(h_b);
free(h_c);
free(h_result);
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
cudaFree(d_states);
return 0;
}
（2）编译
nvcc -o cuda_hello cuda_hello.cu
5. 运行
执行 profile：
nsys profile --stats=true -o my_cuda_profile --trace=cuda,osrt,nvtx --gpu-metrics-devices=0 --cudabacktrace=all --sample=system-wide --force-overwrite true --gpu-metrics-frequency=20000 ./cuda_hello
参数解释：
nsys profile \
--stats=true \ # 生成统计报告
-o my_cuda_profile \ # 输出文件名
--trace=cuda,osrt,nvtx \ # 捕获 CUDA、系统运行时、NVTX 标记
--gpu-metrics-devices=0 \ # 收集 0 号GPU 硬件指标
--cudabacktrace=all \ # 启用 CUDA 调用堆栈跟踪
--sample=system-wide \ # 系统级 CPU 采样
--gpu-metrics-frequency=20000 \ # GPU metrics 采样频率，表示 1秒采样次数，太高可能会导致 Missing Data，太小也可能采集不到
--force-overwrite true \ # 结果文件覆盖写入
./cuda_hello
profile 结果示例详情
root@6b2d3a319f66:/data/cuda_tmp/cuda_helloworld# nsys profile \
--stats=true \
-o my_cuda_profile \
--trace=cuda,osrt,nvtx \
--gpu-metrics-devices=all \
--cudabacktrace=all \
--sample=system-wide \
./cuda_hello
GPU 0: General Metrics for NVIDIA AD10x (any frequency)
GPU 1: General Metrics for NVIDIA AD10x (any frequency)
GPU 2: General Metrics for NVIDIA AD10x (any frequency)
GPU 3: General Metrics for NVIDIA AD10x (any frequency)
GPU 4: General Metrics for NVIDIA AD10x (any frequency)
GPU 5: General Metrics for NVIDIA AD10x (any frequency)
GPU 6: General Metrics for NVIDIA AD10x (any frequency)
GPU 7: General Metrics for NVIDIA AD10x (any frequency)
Collecting data...
Using device 0: NVIDIA L20
Vector add checksum: 1099511627776.00
Matrix multiply checksum: 64.71
Generating '/tmp/nsys-report-6c9b.qdstrm'
[1/8] [========================100%] my_cuda_profile.nsys-rep
[2/8] [========================100%] my_cuda_profile.sqlite
[3/8] Executing 'nvtx_sum' stats report
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Style Range
-------- --------------- --------- ------------ ------------ ---------- ---------- ----------- ------- -------------------------
83.7 69,990,913 1 69,990,913.0 69,990,913.0 69,990,913 69,990,913 0.0 PushPop :Device Memory Allocation
8.2 6,857,564 1 6,857,564.0 6,857,564.0 6,857,564 6,857,564 0.0 PushPop :Vector Add Kernel
4.3 3,609,573 1 3,609,573.0 3,609,573.0 3,609,573 3,609,573 0.0 PushPop :Checksum Calculation
1.0 820,169 1 820,169.0 820,169.0 820,169 820,169 0.0 PushPop :Matrix Memory Operations
0.7 600,315 1 600,315.0 600,315.0 600,315 600,315 0.0 PushPop :Matrix Multiply Sync
0.7 597,025 1 597,025.0 597,025.0 597,025 597,025 0.0 PushPop :D2H Transfer
0.6 541,167 1 541,167.0 541,167.0 541,167 541,167 0.0 PushPop :D2H Transfer (Matrix)
0.6 476,558 1 476,558.0 476,558.0 476,558 476,558 0.0 PushPop :H2D Transfer
0.1 41,909 1 41,909.0 41,909.0 41,909 41,909 0.0 PushPop :Init Random States
0.0 29,049 1 29,049.0 29,049.0 29,049 29,049 0.0 PushPop :Generate Random Matrices
0.0 26,799 1 26,799.0 26,799.0 26,799 26,799 0.0 PushPop :Matrix Multiply Kernel
0.0 12,390 1 12,390.0 12,390.0 12,390 12,390 0.0 PushPop :Host Memory Allocation
0.0 7,920 1 7,920.0 7,920.0 7,920 7,920 0.0 PushPop :Host Result Allocation
[4/8] Executing 'osrt_sum' stats report
Time (%) Total Time (ns) Num Calls Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
-------- --------------- --------- ------------ ----------- -------- ---------- ------------ ----------------------
63.4 138,936,900 11 12,630,627.3 1,124,403.0 2,700 75,039,589 23,446,351.4 poll
34.8 76,361,881 1,578 48,391.6 21,469.5 1,000 17,184,953 483,219.8 ioctl
0.4 934,311 31 30,139.1 7,780.0 3,690 546,276 96,514.7 mmap64
0.4 774,460 9 86,051.1 40,959.0 27,890 450,268 136,998.5 sem_timedwait
0.3 610,639 105 5,815.6 4,370.0 1,800 71,948 8,275.0 open64
0.3 601,237 4 150,309.3 34,719.5 6,480 525,318 250,573.6 fread
0.2 384,709 96 4,007.4 1,625.0 1,000 64,328 7,429.6 fopen
0.1 158,676 1 158,676.0 158,676.0 158,676 158,676 0.0 pthread_cond_wait
0.1 140,376 2 70,188.0 70,188.0 61,459 78,917 12,344.7 pthread_create
0.0 83,518 13 6,424.5 2,130.0 1,210 35,209 9,402.7 mmap
0.0 66,750 10 6,675.0 6,300.0 4,080 11,330 2,373.7 write
0.0 42,659 2 21,329.5 21,329.5 1,530 41,129 28,000.7 fgets
0.0 30,948 1 30,948.0 30,948.0 30,948 30,948 0.0 connect
0.0 30,470 15 2,031.3 1,290.0 1,030 7,090 1,679.8 fclose
0.0 26,670 5 5,334.0 4,370.0 1,850 11,430 3,596.3 open
0.0 21,800 8 2,725.0 1,675.0 1,010 7,810 2,369.5 close
0.0 21,520 2 10,760.0 10,760.0 5,110 16,410 7,990.3 socket
0.0 16,990 6 2,831.7 2,645.0 2,050 3,850 789.4 munmap
0.0 16,150 6 2,691.7 2,315.0 1,650 5,140 1,237.2 read
0.0 13,929 1 13,929.0 13,929.0 13,929 13,929 0.0 bind
0.0 13,920 2 6,960.0 6,960.0 1,280 12,640 8,032.7 pthread_cond_broadcast
0.0 13,150 3 4,383.3 5,360.0 1,790 6,000 2,268.6 pipe2
0.0 6,440 3 2,146.7 1,650.0 1,270 3,520 1,204.4 fcntl
0.0 3,969 1 3,969.0 3,969.0 3,969 3,969 0.0 fwrite
0.0 1,090 1 1,090.0 1,090.0 1,090 1,090 0.0 fopen64
[5/8] Executing 'cuda_api_sum' stats report
Time (%) Total Time (ns) Num Calls Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
-------- --------------- --------- ------------ ----------- -------- ---------- ------------ ---------------------------------
42.3 70,215,465 7 10,030,780.7 129,837.0 104,597 69,478,165 26,213,835.7 cudaMalloc
42.3 70,209,706 7 10,029,958.0 129,547.0 104,317 69,474,486 26,212,576.2 cudaMalloc
4.2 6,897,722 5 1,379,544.4 19,219.0 3,820 6,828,685 3,046,179.7 cudaLaunchKernel
4.2 6,896,842 5 1,379,368.4 19,079.0 3,690 6,828,425 3,046,132.6 cudaLaunchKernel
1.9 3,089,535 2 1,544,767.5 1,544,767.5 594,155 2,495,380 1,344,369.1 cudaDeviceSynchronize
1.9 3,089,335 2 1,544,667.5 1,544,667.5 594,065 2,495,270 1,344,354.9 cudaDeviceSynchronize
1.0 1,583,141 4 395,785.3 382,960.5 226,864 590,356 193,635.6 cudaMemcpy
1.0 1,582,081 4 395,520.3 382,655.5 226,684 590,086 193,540.7 cudaMemcpy
0.6 1,025,114 7 146,444.9 131,787.0 125,677 186,725 23,736.9 cudaFree
0.6 1,023,124 7 146,160.6 131,597.0 125,237 186,425 23,748.1 cudaFree
0.2 272,223 1 272,223.0 272,223.0 272,223 272,223 0.0 cudaGetDeviceProperties_v2_v12000
0.0 900 1 900.0 900.0 900 900 0.0 cuModuleGetLoadingMode
[6/8] Executing 'cuda_gpu_kern_sum' stats report
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
-------- --------------- --------- ----------- ----------- --------- --------- ----------- ----------------------------------------------------------
80.2 2,508,797 1 2,508,797.0 2,508,797.0 2,508,797 2,508,797 0.0 initCurandStates(curandStateXORWOW *, unsigned long long)
19.1 597,855 1 597,855.0 597,855.0 597,855 597,855 0.0 matrixMultiply(const float *, const float *, float *, int)
0.5 15,872 2 7,936.0 7,936.0 7,552 8,320 543.1 generateRandomNumbers(curandStateXORWOW *, float *, int)
0.2 4,864 1 4,864.0 4,864.0 4,864 4,864 0.0 vectorAdd(const float *, const float *, float *, int)
[7/8] Executing 'cuda_gpu_mem_time_sum' stats report
Time (%) Total Time (ns) Count Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Operation
-------- --------------- ----- --------- --------- -------- -------- ----------- ----------------------------
64.3 674,944 2 337,472.0 337,472.0 334,048 340,896 4,842.3 [CUDA memcpy Device-to-Host]
35.7 374,207 2 187,103.5 187,103.5 183,263 190,944 5,431.3 [CUDA memcpy Host-to-Device]
[8/8] Executing 'cuda_gpu_mem_size_sum' stats report
Total (MB) Count Avg (MB) Med (MB) Min (MB) Max (MB) StdDev (MB) Operation
---------- ----- -------- -------- -------- -------- ----------- ----------------------------
8.389 2 4.194 4.194 4.194 4.194 0.000 [CUDA memcpy Device-to-Host]
8.389 2 4.194 4.194 4.194 4.194 0.000 [CUDA memcpy Host-to-Device]
Generated:
/data/cuda_tmp/cuda_helloworld/my_cuda_profile.nsys-rep
/data/cuda_tmp/cuda_helloworld/my_cuda_profile.sqlite
root@6b2d3a319f66:/data/cuda_tmp/cuda_helloworld#
会在控制台输出一系列的统计信息，以及生成 *.nsys-rep 和 *.sqlite 文件。
其中 *.sqlite 文件是性能数据的数据库，nsys stats my_cuda_profile.sqlite 可以提取出一些统计数据，和在控制台里看到的一样。
6. 控制台输出信息
6.1 nvtx_sum 信息
nvtx 信息汇总，是我们在代码里插入的每一个 nvtx range 的耗时汇总显示
6.2 osrt_sum 信息
操作系统系统调用信息汇总
可能和 GPU 有关的函数有：
- poll，CPU 等待 GPU 操作完成
- ioctl，应用程序通过 CUDA API 操作 CUDA 运行时库，最终调用到 cuda 驱动
- mmap64，用于映射 GPU 显存到 CPU 地址空间
6.3 cuda_api_sum 信息
CUDA api 调用耗时统计。大多数 CUDA API（如 cudaMalloc、cudaLaunchKernel、cudaMemcpyAsync）是异步的，意味着 API 调用会立即返回，而实际操作（如显存分配、内核执行、数据传输）会在 GPU 上稍后执行。这里统计的是 API 调用本身的耗时，而非底层操作的实际执行时间。具体来说：
- CPU 端时间：从调用 CUDA API 开始（进入 libcudart.so 库）到 API 返回给用户代码的时间。
- 不包含 GPU 执行时间：例如，cudaLaunchKernel 的耗时仅包括 CPU 提交内核到 GPU 队列的时间，不包括内核在 GPU 上实际运行的时间。
对于异步 API 这里只能看到 CPU 侧的耗时和调用的次数，更深入分析需要结合 NSys 的其他信息。
对于同步 API （譬如：cudaDeviceSynchronize、同步的 cudaMemcpy），这里可以看到 CPU 侧的完整耗时。
6.4 cuda_gpu_kern_sum 信息
CUDA 核函数（GPU 上执行的函数，也叫内核代码）的耗时情况。
6.5 cuda_gpu_mem_time_sum 信息
GPU 内存操作（数据传输）的耗时统计，具体聚焦于 主机（CPU）与设备（GPU）之间的数据拷贝操作。
7. 看 timelines
7.1 常用操作
- Alt+滚轮：左右滑动
- Ctrl+滚轮：放大放小
- Shift + 双击：底部显示 Events View
- Backspace：回退上一步操作
7.2 页面布局
-
左侧有三大块区域：CPU、GPU、Processes，前两个是硬件视角，Processes 是软件视角。
- CPU 区域：显示的是系统中所有物理或逻辑 CPU 核心的利用率和活动。
- GPU 区域：显示的是系统中的 GPU 硬件的活动和指标。
- Processes 区域，则更侧重于 软件（进程和线程）层面 的活动组织。虽然它包含了 CUDA HW（进程触发的 GPU 活动）和 OS runtime libraries 等泳道，但其核心是按应用程序的进程和线程来分类和显示其各自的活动，包括 CPU 端的代码执行、API 调用、以及探查器开销等。它从应用程序的视角来展示 CPU 和 GPU 的交互。
-
软件层面的 CPU 使用情况
- 以 cuda_hello 为例，图中绿色的泳道表示本线程的 CPU 繁忙程度，绿色条越高表示越繁忙，可以注意到和上面的 CPU 硬件统计是一致的，也可以发现本线程被操作系统调度到了多个 CPU 核上执行。
-
GPU 采样和 Processes 事件跟踪：
- Processes 下的 CUDA HW 基于事件跟踪，它只记录明确的开始和结束的事件。如果一个 GPU 活动没有被 Nsys 识别为一个标准的 CUDA 内核或内存传输事件，它就不会显示在这里 Processes 里，但是它可能会被 GPU 采样到，所以某个时刻在 Processes 的 CUDA HW 里是空白，但在 GPU 统计里有颜色。
7.3 GPU metrics
CPU metrics 的信息比较简单，这里介绍下 GPU metrics 里的信息。
GPU Metrics 后面中括号里的 20 KHz 是我们抓取的频率，是我们通过参数 --gpu-metrics-frequency=20000 控制的。
7.3.1 全局硬件时钟频率
（1）GPC Clock Frequency，GPC（Graphics Processing Cluster，图形处理集群）是 NVIDIA GPU 架构中图形渲染核心的顶层组织单元，每个 GPC 包含多个 SM（流式多处理器）、光栅化单元、纹理单元等与图形渲染或者通用计算直接相关的硬件模块，可以说是计算硬件的载体。GPC Clock Frequency 即这些核心计算模块（如 SM、光栅器）的运行时钟频率，单位通常为 MHz 或 GHz。在其他条件都一样的情况下，GPC 频率越高，单位时间内可以处理的指令越多，算得越快，当计算任务充足时，GPC 频率应该达到物理硬件的理论最高值。
（2）SYS（System）Clock Frequency 是 GPU 中辅助性系统模块的运行时钟频率，这些模块不直接参与计算，而是负责 GPU 内部的 “后勤保障” 和外部交互，例如：显存控制器（管理 GPU 与显存的数据传输）；PCIe 接口控制器（负责 GPU 与 CPU / 主板的数据通信）；电源管理模块、温度传感器等辅助电路。对于计算任务来说，这里的频率会影响到数据传输。
7.3.2 计算硬件利用率
LT;DR：我们可以逐级下钻：SM Actives、SM Warp Occupancy、Tensor Active。
（1）GR Active（Graphics Rendering Active，图形渲染活动 ），理论上应该是图形渲染相关的，但由于硬件复用，当计算任务很高的时候，这里的指标也会很高，基本和 SM Actives 同频，我们看 SM Actives 即可。
（2）SM Actives（Streaming Multiprocessor Actives，流多处理器活跃状态）指的是当前处于 “活跃工作状态” 的 SM 单元数量。SM（Streaming Multiprocessor，流多处理器）是 GPU 的核心计算单元，包含 CUDA Core、Tensor Core、缓存、指令调度器等关键组件，是执行计算任务的基本单元。注：SM 内有多个 warp，任意一个 warp 激活了则整个 SM 都算激活，因此我们看的时候还需要进一步下钻到 “SM Warp Occupancy”看 warp 的激活情况。
（3）SM Issue，表示 SM 在每个时钟周期内，warp 指令调度器（Instruction Dispatch Unit）向执行单元发出的指令数量和理论最大值的比例。SM 从活跃的 warp 中选择就绪的指令并发出，让具体的计算单元执行计算。一个 SM 中指令调度器数量因 GPU 架构而异，通常是1~4个或者更多。指令调度器的工作是异步的，发出一条指令之后不需要等它完成，马上可以获取下一条指令。和 warp、计算单元配合可以形成两类并行：
- 计算单元内部实现指令流水线：取指、解码、执行、写回可以在不同的时钟周期内并行。
- warp 之间并行和隐藏延迟：当前 warp 发出的指令如果需要较长时间等待，则调度器可以切换到另一个已经就绪的 warp，并从该 warp 中发出指令，这样可以实现通信与计算的重叠。
（4）Tensor Active：张量核心的活跃度，主要用于深度学习工作负载。在执行大模型推理计算的时候，它通常比较高。
（5）SM Warp Occupancy，SM 里 warp 的占用率。细分为 4 个指标，Vertex/Tess/Geometry Warps in Flight 和 Pixel Warps in Flight 这两是图形相关的，先忽略。
- Compute Warps in Flight (计算 Warp 正在执行)，这个指标表示在 GPU 的 SMs 上，正在执行通用计算 (GPGPU) 任务的活跃 Warp 的比例。这包括所有 CUDA 内核执行的 Warp，无论是进行浮点运算、整数运算、内存操作，还是 Tensor Core 运算（Tensor Active 也会单独显示）。
- Unallocated Warps in Active SMs (活跃 SM 中未分配的 Warp)，指活跃的 SM 中，已占用 SM 资源（如寄存器、共享内存）但未被调度执行任何任务（包括计算、内存操作、同步等）的 warp 数量或比例。这个指标较高说明有 warp 没有执行，计算单元可能处于空闲。
7.3.3 通信硬件利用率
LT;DR：包括几部分：卡内显存带宽、多卡互联带宽、主存和显存带宽
（1）DRAM Bandwidth：显存带宽利用率，即 GPU 卡上显存和 GPU 计算硬件（譬如：SM）之间的数据传输利用率，可以细分为 DRAM Read Bandwidth（读） 和 DRAM Write Bandwidth（写）。
（2）NVLink RX Bandwidth / NVLink TX Bandwidth：NVLink 接收 (RX) 和发送 (TX) 带宽利用率，用于多 GPU 或 GPU 与 CPU 之间的高速互联。
- RX=Receive，TX=Transmit
- NVLink RX Requests Protocol Data: 衡量通过 NVLink 接收的协议数据请求的数量或速率。这些是底层通信协议相关的请求，而不是直接的用户数据。
- NVLink RX Requests User Data: 衡量通过 NVLink 接收的用户数据请求的数量或速率。这直接反映了应用程序数据从发送方到接收方的数据传输请求。
- NVLink RX Responses Protocol Data: 衡量通过 NVLink 接收的协议数据响应的数量或速率。
- NVLink RX Responses User Data: 衡量通过 NVLink 接收的用户数据响应的数量或速率。
（3）PCIe Bandwidth：PCIe 总线带宽利用率（用于 CPU 和 GPU 之间的数据传输），主要用于分析 CPU 内存与 GPU 显存之间的数据传输性能（例如 cudaMemcpy 操作）。
- PCIe RX Throughput (PCIe 接收吞吐量)，即 CPU 到 GPU，对应 H2D，譬如：cudaMemcpyHostToDevice
- PCIe TX Throughput (PCIe 发送吞吐量)，即 D2H，譬如：cudaMemcpyDeviceToHost
7.4 Processes
上图表示 Nsight Systems 监控了 13 个进程，我们重点关注的时候 ./cuda_hello 进程。
7.4.1 CUDA HW
CUDA HW (0000:03:00.0 - NVIDI) 这指的是与该进程关联的 CUDA 硬件。0000:03:00.0 是一个PCIe总线地址，用于标识特定的GPU设备。NVIDI 表明这是一个NVIDIA GPU。
-
55.5% Kernels 这表示在整个CUDA硬件活动中，有 55.5% 的时间花在了 CUDA 核函数（kernels）的执行上。核函数是 GPU 上实际并行执行的计算任务。
- 73.3% initCurandStates 这是在一个或多个 CUDA 核函数中，
initCurandStates函数所占用的时间比例，高达 73.3%。initCurandStates通常与CUDA的随机数生成库cuRAND相关，用于初始化随机数生成器的状态。这表明随机数生成可能是一个主要的计算开销。 - 25.5% matrixMultiply 这是另一个CUDA核函数
matrixMultiply所占用的时间比例，为25.5%。 - 2 kernel groups hid 这表示还有2组核函数的信息被隐藏了，可以通过点击旁边的加号（+）展开查看。展开之后可以看到是一些耗时比较低的操作，如：vectorAdd。
- 73.3% initCurandStates 这是在一个或多个 CUDA 核函数中，
-
44.5% Memory 这表示在整 个CUDA 硬件活动中，有 44.5% 的时间花在了内存操作上。
- 17.7% HtoD memcpy 这指的是主机（Host）到设备（Device）的内存拷贝操作所占用的时间比例，为 17.7%。
HtoD表示数据从 CPU 内存传输到 GPU 显存。 - 82.3% DtoH memcpy 这指的是设备（Device）到主机（Host）的内存拷贝操作所占用的时间比例，为 82.3%。
DtoH表示数据从 GPU 显存传输到 CPU 内存。
- 17.7% HtoD memcpy 这指的是主机（Host）到设备（Device）的内存拷贝操作所占用的时间比例，为 17.7%。
-
NVTX NVTX (NVIDIA Tools Extension) 是一个库，允许开发人员在他们的应用程序中插入自定义事件和范围标记，这些标记可以在 NSys 上展示。这里是 GPU 侧的时间统计，对应的还有 CPU 侧的统计。
7.4.2 Threads
Threads (9) 这表示该进程当前有 9 个活跃的线程，这部分都是 CPU 侧的统计。
-
[572031] cuda_hello 这再次强调了分析的是
cuda_hello进程的主线程或其他相关线程。- OS runtime libraries 这表示操作系统运行时库（如系统调用、线程管理等）所占用的时间。这些是程序运行所必需的底层系统服务。
- NVTX 这里再次出现 NVTX，这是 CPU 侧的时间统计。
- CUDA API 这表示 CUDA API 调用所占用的时间。CUDA API 是应用程序与 GPU 进行交互的接口，这包括启动核函数、内存分配、数据传输等等。
7.5 小块操作的时序图
当我们点击 Processes 下的圆角小块时，会出现一个透明层，同时还会出现很多竖线和白色小三角，譬如，点击 NVTX 里的 Matrix Memory Operations 长条：
-
蓝色竖线和白色小三角：通常表示一个同步点，或者是显示的依赖。
- 竖线，表示每一个操作的开始；
- 白色小三角，表示一个阻塞性操作或者是等待时间，意味着在这个操作之前，需要等待它所依赖的操作完成才能执行。
-
只有蓝色竖线（无白色小三角）：没有小三角意味着这个操作被提交到 GPU 后，CPU 可以里返回并执行后续代码。注：同一个 CUDA 流里的操作都是顺序执行的，即便没有小三角。
-
总结：带小三角的竖线强调的是同步和阻塞，不带小三角的竖线强调的是异步和非阻塞。
7.6 无法识别的核函数
有时候会看到 kernel2 这样的核函数，是 nsys 无法识别导致的，这些函数有可能是 JIT 即时编译生成的。
8. 分析案例
8.1 CPU 和 GPU 的异步
（1）CPU 侧只是将计算或者传输任务提交到 GPU，如要确保执行完成，还需要加上 cudaDeviceSynchronize 函数。
可以看到多个小三角，这些小三角描述了“Matrix Multiply Kernel”这个操作的依赖关系，CPU 侧提交任务的时间很短，真正计算的时间很长。
（2）CPU 开始的时刻和 GPU 真的开始执行可能差很远。
generateRandomNumbers 函数在 CPU 侧提交后，因为它依赖 initCurandStates 初始化的 d_states 数据，所以等到初始化完成才开始执行。
8.2 多级 NVTX
我们可以给 NVTX 增加层级。效果：
示例代码：
示例代码
#include <cuda_runtime.h>
#include <curand.h>
#include <curand_kernel.h>
#include <iostream>
#include <nvtx3/nvToolsExt.h>
#include <stdio.h>
#include <stdlib.h>
// CUDA核函数：向量加法
__global__ void vectorAdd(const float *a, const float *b, float *c, int n)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n)
{
c[idx] = a[idx] + b[idx];
}
}
// CUDA核函数：矩阵乘法
__global__ void matrixMultiply(const float *a, const float *b, float *c, int width)
{
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < width && col < width)
{
float sum = 0.0f;
for (int k = 0; k < width; k++)
{
sum += a[row * width + k] * b[k * width + col];
}
c[row * width + col] = sum;
}
}
// CUDA核函数：初始化随机数
__global__ void initCurandStates(curandState *states, unsigned long long seed)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
curand_init(seed, idx, 0, &states[idx]);
}
// CUDA核函数：生成随机数
__global__ void generateRandomNumbers(curandState *states, float *data, int n)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n)
{
data[idx] = curand_uniform(&states[idx]);
}
}
int main(int argc, char **argv)
{
nvtxRangePushA("Add Test");
nvtxRangePushA("initialize");
const int VECTOR_SIZE = 1 << 20; // 1M elements
const int MATRIX_WIDTH = 1024;
const int MATRIX_SIZE = MATRIX_WIDTH * MATRIX_WIDTH;
float *h_a, *h_b, *h_c; // 主机端数据
float *d_a, *d_b, *d_c; // 设备端数据
curandState *d_states; // 随机数状态
int deviceId;
cudaGetDevice(&deviceId);
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, deviceId);
printf("Using device %d: %s\n", deviceId, deviceProp.name);
nvtxRangePop();
// Host memory allocation
nvtxRangePushA("Host Memory Allocation");
h_a = (float *)malloc(VECTOR_SIZE * sizeof(float));
h_b = (float *)malloc(VECTOR_SIZE * sizeof(float));
h_c = (float *)malloc(VECTOR_SIZE * sizeof(float));
nvtxRangePop();
// 初始化主机端数据
nvtxRangePushA("Init Host Data");
for (int i = 0; i < VECTOR_SIZE; i++)
{
h_a[i] = (float)i;
h_b[i] = (float)(VECTOR_SIZE - i);
}
nvtxRangePop();
// Device memory allocation
nvtxRangePushA("Device Memory Allocation");
cudaMalloc((void **)&d_a, VECTOR_SIZE * sizeof(float));
cudaMalloc((void **)&d_b, VECTOR_SIZE * sizeof(float));
cudaMalloc((void **)&d_c, VECTOR_SIZE * sizeof(float));
nvtxRangePop();
// Host to Device transfer
nvtxRangePushA("H2D Transfer");
cudaMemcpy(d_a, h_a, VECTOR_SIZE * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, VECTOR_SIZE * sizeof(float), cudaMemcpyHostToDevice);
nvtxRangePop();
// 设置向量加法的线程块和网格维度
int threadsPerBlock = 256;
int blocksPerGrid = (VECTOR_SIZE + threadsPerBlock - 1) / threadsPerBlock;
// Vector Add Kernel
nvtxRangePushA("Vector Add Kernel");
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, VECTOR_SIZE);
nvtxRangePop();
// Device to Host transfer
nvtxRangePushA("D2H Transfer");
cudaMemcpy(h_c, d_c, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost);
nvtxRangePop();
// 验证结果
float sum = 0.0f;
for (int i = 0; i < VECTOR_SIZE; i++)
{
sum += h_c[i];
}
printf("Vector add checksum: %.2f\n", sum);
nvtxRangePop();
// ------矩阵乘法部分------
// Matrix memory operations
nvtxRangePushA("Matrix Test");
nvtxRangePushA("Matrix Memory Operations");
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
cudaMalloc((void **)&d_a, MATRIX_SIZE * sizeof(float));
cudaMalloc((void **)&d_b, MATRIX_SIZE * sizeof(float));
cudaMalloc((void **)&d_c, MATRIX_SIZE * sizeof(float));
nvtxRangePop();
// 分配随机数状态内存
cudaMalloc((void **)&d_states, MATRIX_SIZE * sizeof(curandState));
// 初始化随机数生成器
dim3 threads(16, 16);
dim3 blocks(MATRIX_WIDTH / threads.x, MATRIX_WIDTH / threads.y);
// Initialize random states
nvtxRangePushA("Init Random States");
initCurandStates<<<blocks, threads>>>(d_states, 1234ULL);
nvtxRangePop();
// Generate random matrices
nvtxRangePushA("Generate Random Matrices");
generateRandomNumbers<<<blocks, threads>>>(d_states, d_a, MATRIX_SIZE);
generateRandomNumbers<<<blocks, threads>>>(d_states, d_b, MATRIX_SIZE);
nvtxRangePop();
// 同步以确保随机数生成完成
cudaDeviceSynchronize();
// 调用矩阵乘法核函数
nvtxRangePushA("Matrix Multiply Kernel");
matrixMultiply<<<blocks, threads>>>(d_a, d_b, d_c, MATRIX_WIDTH);
nvtxRangePop();
// 同步以确保矩阵乘法完成
nvtxRangePushA("Matrix Multiply Sync");
cudaDeviceSynchronize();
nvtxRangePop();
// 分配主机端内存用于结果矩阵
nvtxRangePushA("Host Result Allocation");
float *h_result = (float *)malloc(MATRIX_SIZE * sizeof(float));
nvtxRangePop();
// 设备到主机的数据传输
nvtxRangePushA("D2H Transfer (Matrix)");
cudaMemcpy(h_result, d_c, MATRIX_SIZE * sizeof(float), cudaMemcpyDeviceToHost);
nvtxRangePop();
// 计算结果校验和
nvtxRangePushA("Checksum Calculation");
float checksum = 0.0f;
for (int i = 0; i < MATRIX_SIZE; i++)
{
checksum += h_result[i];
}
printf("Matrix multiply checksum: %.2f\n", checksum);
nvtxRangePop();
nvtxRangePop();
nvtxRangePushA("free memory");
// 释放内存
free(h_a);
free(h_b);
free(h_c);
free(h_result);
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
cudaFree(d_states);
nvtxRangePop();
return 0;
}
9. 抓取部分 iteration
注：iteration 是指生成一个 token 的全过程。
在做大模型推理性能分析时，过多的数据会影响分析的效率，我们通常只抓取部分 iteration 做分析。可以通过使用 cudaProfilerStart 和 cudaProfilerStop 接口实现，下面以 python 为例做介绍。
（1）在代码中根据条件调用 cudaProfilerStart 和 cudaProfilerStop 函数，譬如：
def event_loop_normal(self):
self._profile_iter = 0
while True:
...
self._profile_iter += 1
if self._profile_iter == 1000:
torch.cuda.cudart().cudaProfilerStart()
elif self._profile_iter == 1050:
torch.cuda.cudart().cudaProfilerStop()
...
我们也可以将上述例子代码中的 1000 和 1050 设计为参数或者环境变量，实现更灵活的配置。
（2）以 nsys profile 启动程序，譬如：
nsys profile \
--gpu-metrics-frequency=10000 \
--gpu-metrics-devices=0 \
--force-overwrite true \
-o sglang_profile_h20_tp1_overlap_stage \
--capture-range=cudaProfilerApi \
--capture-range-end=stop \
python3 -m sglang.launch_server \
--model-path /data/xxxx \
--port 30000 &> tp1_overlap.log
上述 shell 命令里 --capture-range-end=stop 如果不加上会导致进程完成捕获之后就退出，如果加上了，则表示当调用 cudaProfilerStop 的时候，就停止捕获，虽然进程不退出，但进程的整个生命周期只捕获一次，如果在应用程序的生命周期里要捕获多个 iteration 区间，则需要使用 repeat，repeat 可以指定捕获 N 次，或者无限次，每一次都是 cudaProfilerStart 和 cudaProfilerStop 配对调用。譬如：
PROFILE_START_STOP=10-20,30-40,66 nsys profile \
--gpu-metrics-frequency=10000 \
--gpu-metrics-devices=0 \
--force-overwrite true \
-o sglang_profile_h20_tp1_normal_stage \
--capture-range=cudaProfilerApi \
--capture-range-end=repeat:3 \
python3 -m sglang.launch_server \
--model-path /data/7b/checkpoint-743 \
--disable-overlap-schedule \
--port 30000 &> tp1_normal.log
其中 PROFILE_START_STOP 是通过环境变量实现灵活的配置抓取 iteration 区间。--capture-range-end=repeat:3 表示抓取 3 个 range，每一个 range 都会有单独的 rep 文件。如果指定为 --capture-range-end=repeat:0，则需要进程退出时才能生成 rep 文件。
浙公网安备 33010602011771号