cuda学习笔记：运行分析1.0

最新推荐文章于 2026-03-20 15:07:24 发布

原创最新推荐文章于 2026-03-20 15:07:24 发布 · 1.2k 阅读

本内容遵循CC 4.0 BY-SA版权协议

一、分析工具：Nsight

======== Warning: nvprof is not supported on devices with compute capability 8.0 and higher.
                  Use NVIDIA Nsight Systems for GPU tracing and CPU sampling and NVIDIA Nsight Compute for GPU profiling.
                  Refer https://developer.nvidia.com/tools-overview for more details.

原来如此，现在已经改成nsight工具来分析使用了，nsight分为systems和compute（话说installer居然不给我配全环境变量，有点懵，大家可以在终端试试下面两个命令，测试这两个工具是否能运行）

nsys -v
// NVIDIA (R) Nsight Compute Command Line Profiler
// Copyright (c) 2018-2025 NVIDIA Corporation
// Version 2025.2.1.0 (build 35987062) (public-release)
ncu -v
// NVIDIA Nsight Systems version 2025.1.3.140-251335620677v0

Nsight Systems：系统及分析器，捕捉整个应用从主机到GPU的完整执行轨迹，包括CPU调度，cuda调用，线程同步等

功能：

1.命令采集代理：将应用各种事件写入时间线记录，比如os线程调度，cudaapi，nvtx标记等

2.目标茶庄代理：在被分析进程内部插桩，收集GPU活动：kernel launch、memcpy、硬件事件计数器等

3.可视化：生成.qdrop或者nsys-rep文件，后续可以用nsys-ui查看

等等

Nsight Compute：聚焦cuda内核的深度剖析

1.深入剖析单个或一组cudakernel 的硬件利用情况，通过丰富的指标帮助你发现内核级性能瓶颈。

2.分析寄存器分配、SM occupancy、内存访问带宽、缓存命中率、指令吞吐、warp 发散等

这里随便写了个三个核函数的代码，让nsys监控运行

#include <stdio.h>

__global__ void mul(float* a,float value){
    a[threadIdx.x] = value;
    a[threadIdx.x] *=1.5;
    
}
__global__ void add(float* a,float value){
    a[threadIdx.x] = value;
    a[threadIdx.x] +=1;

}
__global__ void add_ab(float* a,float* b,float* c){ 
    c[threadIdx.x] = a[threadIdx.x] + b[threadIdx.x];
}

void print(float *c,int nElem){
    for (int i = 0; i < nElem; i++){
        printf("%f ",c[i]);
    }
}

int main(void){
    cudaSetDevice(0);
    float* a,*b ,*c_d,*c;
    int nElem = 100;
    size_t sz = nElem*sizeof(float);
    cudaMalloc((void**)&a,sz);
    cudaMalloc((void**)&b,sz);
    cudaMalloc((void**)&c_d,sz);
    c = (float*)malloc(sz);
    cudaMemset(a,0,sz);
    cudaMemset(b,0,sz);
    cudaMemset(c_d,0,sz);
    add<<<1,nElem>>>(a,1.0f);
    mul<<<1,nElem>>>(b,1.0f);
    cudaDeviceSynchronize();
    add_ab<<<1,nElem>>>(a,b,c_d);
    cudaDeviceSynchronize();
    cudaMemcpy(c,c_d,sz,cudaMemcpyDeviceToHost);
    print(c,nElem);
    cudaDeviceReset();
    getchar();
    return 0;
}

注意需要在管理员身份启动的终端才能完全运行nsys监控

nsys profile --output=report --trace=cuda  --gpu-metrics-device=0  ./hello.exe

生成了report.nsys-rep文件，用nsys-ui report.nsys-rep打开它：下图

然后用ncu分析一下：

有点多，只截取了一部分，可以看到给出的信息非常详细

 add_ab(float *, float *, float *) (1, 1, 1)x(100, 1, 1), Context 1, Stream 7, Device 0, CC 8.9
    Section: GPU Speed Of Light Throughput
    ----------------------- ----------- ------------
    Metric Name             Metric Unit Metric Value
    ----------------------- ----------- ------------
    DRAM Frequency                  Ghz         7.69
    SM Frequency                    Ghz         2.21
    Elapsed Cycles                cycle        3,678
    Memory Throughput                 %         0.84
    DRAM Throughput                   %         0.84
    Duration                         us         1.66
    L1/TEX Cache Throughput           %        17.15
    L2 Cache Throughput               %         0.47
    SM Active Cycles              cycle        69.96
    Compute (SM) Throughput           %         0.04
    ----------------------- ----------- ------------

    OPT   This kernel grid is too small to fill the available resources on this device, resulting in only 0.0 full
          waves across all SMs. Look at Launch Statistics for more details.

    Section: GPU Speed Of Light Roofline Chart
    INF   The ratio of peak float (FP32) to double (FP64) performance on this device is 64:1. The workload achieved
          close to 0% of this device's FP32 peak performance and 0% of its FP64 peak performance. See the Profiling
          Guide (https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html#roofline) for more details on
          roofline analysis.

ncu会具体监控到每个核函数的运行记录，便于我们发现每个函数的瓶颈所在。

二、cuda程序运行流程

程序流程

1.执行文件加载（exe文件和相关库dll文件）与初始化，在代码中第一次调用任何cuda代码，比如cudaSetDevice时，libcudart会自动调用一个Driver api：cuInit（）初始化cuda driver，并在GPU上创建一个cuda context，负责维护程序的stream，events等，类似于在GPU上创建一个虚拟进程

libcudart 是 NVIDIA CUDA（计算统一设备架构）平台的重要组成部分。它代表 CUDA 运行时库，提供一组管理和执行 CUDA 程序所需的运行时功能和工具。这些功能包括 GPU 内存管理、设备管理、流管理和错误处理。libcudart 依赖于另一个名为 libcuda 的库。libcudart 提供高级运行时支持，而 libcuda 是 NVIDIA CUDA 驱动程序的一部分，并提供用于与 GPU 通信的低级接口。本质上，libcuda 管理 GPU 的初始化、CUDA 程序的执行和内存访问。

2.host端提交指令。cuda的命令大致分为两类，一类是准备数据的内存操作，另一类是核函数，都通多cuda stream提交。

当调用内存分配时，比如cudaMalloc，runtime会调用cuMemAlloc()，driver在context中在GPU的内存空间分配一块虚拟地址，并记录映射

调用内存拷贝时，如果是同步拷贝，cudaMemcpy -> 调用cuMemcpyXtoX(),host线程会阻塞知道GPU的DMA引擎把数据从主机内存拷贝到设备内存完成（pinned）

如果是异步拷贝调用cuMemcpyXtoXAsyc()，driver会把dma发送命令放入stream的命令队列（Command BUffer），然后立即返回，具体的拷贝操作由GPU dma引擎后台执行

当调用kernel launch时，核函数调用，kernelfunc<<<g,b>>>(args) -> cuLaunchKernel(function_handle,g.x,g.y,g.z,b.x,b.y,b.z,kernal_params)。Driver会把launch kernel命令编码进当前stream的command buffer里（一系列二进制命令，存在 Host 端的一块缓冲区）

Stream Flush和Submission等命令，当command buffer达到一定大小或需要同步操作，driver会把buffer交给gpu。

3.device执行

1. 内核命令排队：gpu的驱动内核模块nvidia.ko/nvlddmkm.sys接受到command buffer后，将他放入context的submission queue提交队列中

2.CP（command processor），它会不断通submission queue取命令，解析成对Kernel Scheduler、DMA Engine、Memory Controller的操作。对于 Kernel Launch，CP 会把需要调度的线程块（thread-block）信息发给Kernel Scheduler

3.Grid调度，kernel scheduler收到一个kernel grid后，根据sm的可用资源，寄存器，共享内存等，把thread-block划分到各个sm上执行，并在每个sm的warp scheduler里排队

4.在sm上执行，每个sm有多个warp scheduler，负责每个cycle发射一个warp（32个线程）的指令到cuda core里（在alu，fpu，tensor core等硬件上执行）。因为是simt模型来管理，当发生分化时，蛋哥warp内部lane会静默（masked-off）

5.内存访问，当kernel里有全局内存访问（ld.global、st.global），指令会通过 L1→L2Cache→DRAM Controller 路径访问显存，访问会被硬件 coalesce，并发给 DRAM，通过 Memory Partition Units并发读写。

4.内存拷贝回host并同步、

1.内存拷贝回传，调用cudaMemcpyAsync（DtoH）时，同样在stream里enqueue一个dma命令，CP解析后由Async Copy Engine（DMA Engine）把数据通过 PCIe 拷回主机内存。

2.事件同步，cudaEventRecord（）会在stream里出入一个事件命令，等待前面操作执行完再在gpu写一个标志cudaEventSynchronize（）则会阻塞host，查询该事件的完成状态（driver 内部 poll或者Query GPU register），直到完成返回

3.设备同步，各种sync命令执行完各个设备上的命令

5.退出

1、Context 销毁，当进程调用cudaDeviceReset（）或者正常退出时，CUDA Runtime 会调用cuCtxDestroy（）销毁context。driver会撤销所有未完成的命令、释放gpu上分配的内存、卸载 module。

2.驱动与硬件清理，GPU 上的Submission Queue、command processor 状态被清空。驱动内核模块释放对GPU的引用，若无其他进程使用，则最终设备回到空闲状态。

AI总结：

用户进程 (CUDA Runtime)
├─> cuInit() / cuDeviceGet() / cuCtxCreate()
├─> 模块加载 (PTX → JIT → SASS)
├─> cudaMalloc / cuMemAlloc()
├─> cudaMemcpyAsync / cuMemcpyHtoDAsync() ─┐
├─> myKernel<<<>>> → cuLaunchKernel() ├─> Command Buffer (Host)
└─> cudaStreamSynchronize() └─[ ioctl / DeviceIoControl ]

内核驱动 (nvidia.ko/nvlddmkm.sys)
└─> Submission Queue

GPU 硬件
├─> Command Processor 解码命令
├─> Kernel Scheduler 分发到各 SM
├─> Warp Scheduler → CUDA Cores 执行
├─> Async Copy Engine 执行 DMA
└─> 缓存层次 + DRAM Controller 处理内存访问

Host ↔ GPU 交互通过 PCIe (或 NVLink)，指令与数据均封装在 DMA Transaction 和命令队列里。

从用户代码的 <<<>>> 到最终在每个 SM 上发射 warp、再到 memcpy 回主机，这整条链路里涉及 Runtime→Driver API→内核模块→GPU Command Processor→SM 的全过程。

Driver和Runtime

下面在介绍一下Driver和Runtime

cuda分两层编程接口，运行时的Runtime api和驱动层的driver api

Runtime Api：

以cuda开头，会在内部调用driver层的api，他就是我们正常编码所使用的接口，包括我们运行runtime函数时自动创建context，通过__global__和nvcc编译和<<<>>>launch kernel

Driver Api：

以cu开头，是低级的接口，需要显式管理程序，包括上下文的创建，cuCtxCreate/Destroy()等，但是这样可以做到更细粒度的控制，可以选择性的管理stream event等

二者可以混合使用，但是Runtime适合入门

三、程序优化1.0

#include <stdio.h>


__global__ void sumMatrix2D(float *a,float*b,float* c,int nx,int ny){
    unsigned int ix = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int iy = blockIdx.y * blockDim.y + threadIdx.y;
    // 计算在数据在内存中的索引 
    unsigned int idx = iy * nx + ix;
    if (ix < nx && iy < ny){
        c[idx] = a[idx] + b[idx];
    }


}

int main(int argc,char **argv){
    cudaSetDevice(0);

    int nx = 1<<14;
    int ny = 1<<14;
    int nxy = nx * ny;

    float *a_d,*b_d,*c_d,*c_h;

    c_h = (float*)malloc(nxy*sizeof(float));
    cudaMalloc((void**)&a_d,nx*ny*sizeof(float));
    cudaMalloc((void**)&b_d,nx*ny*sizeof(float));
    cudaMalloc((void**)&c_d,nx*ny*sizeof(float));

    cudaMemset(a_d,0,nx*ny*sizeof(float));
    cudaMemset(b_d,0,nx*ny*sizeof(float));
    cudaMemset(c_d,0,nx*ny*sizeof(float));

    int dimx = 16,dimy=16;
    if (argc > 1){
        dimx = atoi(argv[1]);
        dimy = atoi(argv[2]);
    }

    dim3 block(dimx,dimy);
    dim3 grid((nx+block.x-1)/block.x,(ny+block.y-1)/block.y);
    sumMatrix2D<<<grid,block>>>(a_d,b_d,c_d,nx,ny);
    cudaDeviceSynchronize();
    cudaMemcpy(c_h,c_d,nxy*sizeof(float),cudaMemcpyDeviceToHost);

    cudaFree(a_d);
    cudaFree(b_d);
    cudaFree(c_d);
    free(c_h);
    return 0;
}

这里给出一个矩阵加法的cuda程序

这是16 * 16的块的ncu检测：

  sumMatrix2D(float *, float *, float *, int, int) (1024, 1024, 1)x(16, 16, 1), Context 1, Stream 7, Device 0, CC 8.9
    Section: GPU Speed Of Light Throughput
    ----------------------- ----------- -------------
    Metric Name             Metric Unit  Metric Value
    ----------------------- ----------- -------------
    DRAM Frequency                  Ghz          7.99
    SM Frequency                    Ghz          2.29
    Elapsed Cycles                cycle    31,746,246
    Memory Throughput                 %         91.57
    DRAM Throughput                   %         91.57
    Duration                         ms         13.83
    L1/TEX Cache Throughput           %         17.49
    L2 Cache Throughput               %         15.16
    SM Active Cycles              cycle 31,637,299.04
    Compute (SM) Throughput           %         10.96
    ----------------------- ----------- -------------

DRAM Frequency 是GPPU全局显存的时钟频率

SM Frequency 是SM的核心时钟频率，表示每个SM的执行单元理论最高launch速率

Elapsed Cycles 表示从kernel启动到结束的总GPU时钟周期数
Memory Throughput 表示内存系统整体吞吐率占理论峰值的百分比
DRAM Throughput 表示物理DRAM不含L2L1cache实际吞吐占峰值的百分比
Duration 表示kernel的执行时长
L1/TEX Cache Throughput 表示L1 及纹理缓存的带宽利用率
L2 Cache Throughput 表示L2 缓存的带宽利用率
SM Active Cycles 表示GPU SM 在活跃（有指令发射）的周期数。几乎等于Elapsed Cycles，说明几乎全程都有内核在跑。
Compute (SM) Throughput SM执行单元在理论最大指令吞吐能力下的利用率。10.96%较低，说明算力资源并未跑满，多半在等待内存。

好的，让我们来测算不同的block大小对应duration：

(1024, 1024, 1)x(16, 16, 1) Duration 13.83

(512, 512, 1)x(32, 32, 1) Duration 13.55

(512, 1024, 1)x(32, 16, 1) Duration 13.98

(1024, 512, 1)x(16, 32, 1) Duration 13.83

emm感觉现在学的知识还不足以分析出差异的原因，没事慢慢学

标签