cuda调试（一）vs2019-windows-Nsight system--nvtx使用，添加nvToolsExt.h文件

cuda调试

由于在编程过程中发现不同的网格块的结构，对最后的代码结果有影响，所以想记录一下解决办法。

CUDA的Context、Stream、Warp、SM、SP、Kernel、Block、Grid

cuda context (上下文)

context类似于CPU进程上下，表示由管理层 Drive 层分配的资源的生命周期，多线程分配调用的GPU资源同属一个context下，通常与CPU的一个进程对应。

CUDA Stream

CUDA Stream是指一堆异步的CUDA操作，他们按照host代码调用的顺序执行在device上。
Stream维护了这些操作的顺序，并在所有预处理完成后允许这些操作进入工作队列，同时也可以对这些操作进行一些查询操作。
这些操作包括host到device的数据传输，launch kernel以及其他的由host发起由device执行的动作。
这些操作的执行总是异步的，CUDA runtime会决定这些操作合适的执行时机。我们则可以使用相应的cuda api来保证所取得结果是在所有操作完成后获得的。同一个stream里的操作有严格的执行顺序，不同的stream则没有此限制。

CUDA API可分为同步和异步两类，同步函数会阻塞host端的线程执行，异步函数会立刻将控制权返还给host从而继续执行之后的动作。

当我们使用CUDA异步函数与多流（Multi Stream）时，多线程间既可以实现并行进行数据传输与计算，如下图所示。不过需要注意的是， CUDA runtime API默认的default stream是同步串行的，且一个进程内的所有线程都在default stream下，需要显式声明default之外的Stream才可以实现多流并发。

显卡硬件架构：SM、SP、Warp

具体到nvidia硬件架构上，有以下两个重要概念：

SP（streaming processor）：最基本的处理单元，也称为CUDA core。最后具体的指令和任务都是在SP上处理的。GPU进行并行计算，也就是很多个SP同时做处理。

SM（streaming multiprocessor）：多个SP加上其他的一些资源组成一个SM，也叫GPU大核，其他资源如：warp scheduler，register，shared memory等。SM可以看做GPU的心脏（对比CPU核心），register和shared memory是SM的稀缺资源。CUDA将这些资源分配给所有驻留在SM中的threads。因此，这些有限的资源就使每个SM中active warps有非常严格的限制，也就限制了并行能力。如下图是一个SM的基本组成，其中每个绿色小块代表一个SP。

每个SM包含的SP数量依据GPU架构而不同，Fermi架构GF100是32个，GF10X是48个，Kepler架构都是192个，Maxwell都是128个。当一个kernel启动后，thread会被分配到很多SM中执行。大量的thread可能会被分配到不同的SM，但是同一个block中的thread必然在同一个SM中并行执行。

Warp调度

一个SP可以执行一个thread，但是实际上并不是所有的thread能够在同一时刻执行。Nvidia把32个threads组成一个warp，warp是调度和运行的基本单元。warp中所有threads并行的执行相同的指令。一个warp需要占用一个SM运行，多个warps需要轮流进入SM。由SM的硬件warp scheduler负责调度。目前每个warp包含32个threads（Nvidia保留修改数量的权利）。所以，一个GPU上resident thread最多只有 SM*warp个。

同一个warp中的thread可以以任意顺序执行，active warps被SM资源限制。当一个warp空闲时，SM就可以调度驻留在该SM中另一个可用warp。在并发的warp之间切换是没什么消耗的，因为硬件资源早就被分配到所有thread和block，所以新调度的warp的状态已经存储在SM中了。

每个SM有一个32位register集合放在register file中，还有固定数量的shared memory，这些资源都被thread瓜分了，由于资源是有限的，所以，如果thread比较多，那么每个thread占用资源就叫少，thread较少，占用资源就较多，这需要根据自己的要求作出一个平衡。

软件架构：Kernel、Grid、Block

上面的context与stream类似进程、线程的概念，具体到我们如何调用GPU上的线程实现我们的算法，则是通过Kernel实现的。在GPU上调用的函数成为CUDA核函数（Kernel function），核函数会被GPU上的多个线程执行。我们可以通过如下方式来定义一个kernel：

func_name<<<grid, block>>>(param1, param2, param3....);

Grid：由一个单独的kernel启动的所有线程组成一个grid，grid中所有线程共享global memory。Grid由很多Block组成，可以是一维二维或三维。

Block：一个grid由许多block组成，block由许多线程组成，同样可以有一维、二维或者三维。block内部的多个线程可以同步（synchronize），可访问共享内存（share memory）。

CUDA中可以创建的网格数量跟GPU的计算能力有关，可创建的Grid、Block和Thread的最大数量如下所示：

以上是引用：

https://zhuanlan.zhihu.com/p/266633373

通过deviceQuery可以知道，每个block中只有1024个thread，而32 × 32× 32 = 32768 > 1024：会出现以下错误，运行时参数传递得太大了，出现这类问题后，cuda仍可继续提供服务，仅单纯拒绝了启动核函数。

cudaErrorInvalidConfiguration = 9，"invalid configuration argument"

所有CUDA kernel的启动都是异步的，当CUDA kernel被调用时，控制权会立即返回给CPU。在分配Grid、Block大小时，我们可以遵循这几点原则：

1. 保证block中thread数目是32的倍数。这是因为同一个block必须在一个SM内，而SM的Warp调度是32个线程一组进行的。
2. 避免block太小：每个blcok最少128或256个thread。
3. 根据kernel需要的资源调整block，多做实验来挖掘最佳配置。
4. 保证block的数目远大于SM的数目。

配置Nsight system

低开销的性能分析工具，Nvidia nsight Systems旨在提供开发人员优化其软件所需的洞察力。在工具中可视化活动数据，以帮助用户调查瓶颈，避免推断误报，并以更高的性能提高概率进行优化。用户将能够识别问题，例如GPU不足、不必要的GPU同步、CPU并行化不足，甚至目标平台的CPU和GPU之间的算法异常。

1. 打开该程序
2. 上面的 command line with arguments中填写 要调试的cuda项目的exe文件我的目录是F:\E_cuda\test1\testCUDA\x64\Debug\testCUDA.exe
   注意这个路径需要是全英文的
   3.下面的 working directory 填写 cuda中的Debug文件，我的cuda中的默认目录是
   C:\ProgramData\NVIDIA Corporation\CUDA Samples\v11.0\bin\win64\Debug

F:\E_cuda\test1\testCUDA\x64\Debug\testCUDA.exe

然后start就行了

但是还不知道怎么去看具体的kernel

这个Report包含5部分内容：

1. Analysis Summary （分析总结，内容非常全面，包含了Target的详细信息，Process summary, Module summary, Thread summary, Environment Variables, CPU info, GPU info等等）
2. Timeline View （展示CPU/GPU各个核的工作时间线，一般用来来勘察模型训练或者推理的瓶颈在哪里）
3. Diagnostics Summary （顾名思义，诊断总结。就是程序在运行中做了什么，有什么warning , error，或者message的，都在这里汇总）
4. Symbol Resolution Logs（暂时不知道是干嘛的）
   Files (执行结果的log 文件：pid_stdout.log，& 执行出错的log 文件pid_stderr.log)

简单看看Timeline view。如下，这里有三个CPU核在工作，它们启动和停止的时间可以从timeline上看到。还可以看到下面有三个Thread的时间线。

https://blog.csdn.net/NXHYD/article/details/112915968

这位博主的图和说明

我在b站上找的视频发现别人用的是 nsight system profiles 这个，（还没找到是啥，只有一个visoul profiler—需要java的环境）但是他讲的时候又说是NVTX，有点晕

https://www.bilibili.com/video/BV13w411o7cu/?spm_id_from=333.337.search-card.all.click&vd_source=0a4d8c47345ce0df71cd9cdb01575134

就是这个
以下是我使用visoul profiler—的报错

英伟达的官方文档关于nsight-system的

https://docs.nvidia.com/nsight-systems/2020.3/profiling/index.html

NVTX是一种CUDA Profiler的工具

可以用于在CUDA程序中进行标记和注释，以便更好地理解和优化程序的性能。以下是使用NVTX的一些示例代码和步骤：

1. 在CUDA代码中包含nvToolsExt.h文件。
   #include <nvToolsExt.h>
   解决：右键cuda项目——属性——配置属性——C/C++——常规——附加包含目录
   可能遇到的问题：无法在 Visual Studio 中打开 nvToolsExt.h 文件，可能是因为 Visual Studio 找不到 CUDA 的 include 目录。

错误示范
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\include
but我这里没有这个头文件

C:\ProgramData\NVIDIA Corporation\CUDA Samples\v11.0\common\inc
这个目录下也没有

正确示范 可以通过在vs项目属性里面进行设置

第一步：在项目属性页，打开vc++目录（c程序是c目录会略微有些不同），打开外部包含目录，输入：$(NVTOOLSEXT_PATH)\include

第二步，打开链接器–常规----附加库目录，输入：$(NVTOOLSEXT_PATH)\lib\$(Platform)

第三步，打开链接器–输入—附加依赖项，输入：nvToolsExt64_1.lib

在 .cu 文件中注释代码，还应该设置：打开CUDA C/C+±-common—附加依赖项，输入：$(NVTOOLSEXT_PATH)\include


英伟达官方NVTX文档

https://docs.nvidia.com/nsight-visual-studio-edition/nvtx/index.html

2. 在需要标记的代码块前后插入标记函数。
   nvtxRangePushA("My Code Block"); // 开始标记代码块 nvtxRangePop(); // 结束标记

3. 可以使用nvtxEventAttributes_t结构体自定义标记的颜色和描述等信息。

nvtxEventAttributes_t eventAttrib = {0};
eventAttrib.version = NVTX_VERSION;
eventAttrib.size = NVTX_EVENT_ATTRIB_STRUCT_SIZE;
eventAttrib.colorType = NVTX_COLOR_ARGB;
eventAttrib.color = 0xFF00FF00;
eventAttrib.messageType = NVTX_MESSAGE_TYPE_ASCII;
eventAttrib.message.ascii = "My Custom Event";
nvtxRangePushEx(&eventAttrib);
// 代码块
nvtxRangePop();

4. 还可以使用nvtxMarkA()函数在代码中插入注释。

nvtxMarkA("My Comment");

5. 在运行程序时，可以使用nvprof或Nvvp等工具来查看标记和注释。
   注意：使用NVTX可能会对程序的性能产生一定的影响

cuda运行时间记录

在python里面
异步计算的结果是没有同步的时间测量是不准确的。要获得精确的测量值，应该在测量前调用 torch.cuda.synchronize()，或者使用 torch.cuda.Event 记录时间，如下所示：

 start_event = torch.cuda.Event(enable_timing=True)
end_event = torch.cuda.Event（enable_timing=True）
start_event.record()# 在这里运行一些东西end_event.record()
torch.cuda.synchronize() # 等待事件被记录！
elapsed_time_ms = start_event.elapsed_time(end_event)

https://blog.csdn.net/dedell/article/details/121574306

先这样，后续会补充，老师让我去搞二维的共享内存的程序，祝我好运！！！