CUDA C

CUDA C 编程
- 名词的意思
- 1.0并行计算与GPU架构
- 1.1 异构计算与CUDA
- 2.1 CUDA编程模型（一）
- 核函数计时
- 组织并行线程
- GPU设备信息
使用Nvidia Profiler
- Nsys
- - NVTX
  - 结合launch.py
- Nsight Compute

CUDA C 编程

最近被安利到一个特别好的博客，看上去非常全面------>谭升的博客
相关书籍推荐---->书籍在blog的尾部

名词的意思

在读GNNAdvisor时，一直对里面的block，thread，warp概念不是很了解，后来才发现这是CUDA里面的一些专有名词。
可以简单参考------>CUDA编程——GPU架构，由sp，sm，thread，block，grid，warp说起
底下的讨论得知，一个SM只执行一个Warp，但是可以加载多个Warp？因为其他Warp可能处于其他状态（挂起，就绪，执行）。一个Warp中threads数量一定（32），一个block中threads数量不一定。
实际上同时执行的threads应该称为resident threads，数量是SM32。 active warp是指已经分配给SM的warp，并且该warp需要的资源（寄存器）也已经分配，这个warp的称为active warp，这个warp通过warp active_mask把此warp中的thread的active bit置1，active thread就是active warp32。

thread：一个CUDA的并行程序会被以许多个threads来执行。
block：数个threads会被群组成一个block，同一个block中的threads可以同步，也可以通过shared memory通信。
grid：多个blocks则会再构成grid。
warp：GPU执行程序时的调度单位，目前cuda的warp的大小为32，同在一个warp的线程，以不同数据资源执行相同的指令,这就是所谓 SIMT。

1.0并行计算与GPU架构

延迟
吞吐量
带宽

1.1 异构计算与CUDA

异构计算
Hello，world–CUDA

2.1 CUDA编程模型（一）

编译cuda文件.cu

nvcc -arch sm_86 xxx.cu -o xxx(类似于exe)
./xxx

但是出现

The program ‘nvcc’ is currently not installed. You can install it by typing:
sudo apt install nvidia-cuda-toolkit

第一想法是nvcc没有在conda环境中装，于是在conda上装一下，参考nvcc_problem，所以是要装一个cudatoolkit-dev，参考git-cudatoolkit-dev,装一下即可。
其实正确的做法应该是找到系统安装的cuda路径，将环境变量配置好，nvcc环境在

/usr/local/cuda-11.1/bin

或者

whereis nvcc

找到nvcc路径，加入即可。

C++文件include其他文件夹头文件，参考#include不同文件夹下头文件，但是需要注意是斜杠的方向。

核函数计时

在使用nvcc编译代码时，图方便就直接

nvcc xxx.cu xxx

结果出现

ERROR: sum_arrays_timer.cu:39,code:222,reason:the provided PTX was compiled with an unsupported toolchain.

搜了一下发现驱动程序和编译链不匹配，也没管直接把-arch 和sm参数加进去发现就可以了，没太懂-------------？
然后在计时和nsys的时候，第一遍按作者将memcpy DtoH 也加入了cpu计时器当中，使用nsys跑出来的结果

再将memcpy HtoD放在计时器外面，nsys结果

可以看到cpu计时器少去的时间正好是CUDA memory统计的时间，而后者的cpu计时与CUDA API统计的cudaLaunchKernel时间基本一致（按博文里说的这里应该会大一点，也非常符合），同时CUDA Memory统计的总时间与CUDA API统计的cudaMemcpy时间也是一致的，唯一有问题的就是CUDA kernel统计的时间，与哪个都对不上-----------------------?
（查文档）

组织并行线程

这里面需要了解

threadsIdx.x,threadIdx.y,：线程所在block的xy坐标
blockIdx.x,blockIdx.y,：block所在grid的xy坐标
blockDim.x,blockDim.y,：grid的维度
block.x,block.y：block的维度，上者在device中使用，这个在host中使用好像。。。

GPU设备信息

很多，有些不太懂，查文档

使用Nvidia Profiler

Nsys

关于nsys的具体用法及参数，可以参考Nsys documents–Nvidia
形式为

nsys [command_switch][optional command_switch_options][application] [optional application_options]

我刚开始用的是

/usr/local/cuda-11.1/bin/nsys profile --stats=true ./sum_arrays

可以自动生成.qdrep文件，该文件需要Nsight Systems来可视化，于是需要安装该软件进行本地查看。
--------->Nsys 可视化Mac
安装该blog一步步安装即可，这里面会出现mac权限的问题，按照教程
‘System Preferences’ → ‘Security & Privacy’ → ‘General’，然后点击允许打开即可。
最后将文件拖进去就可以查看，但是现在还有一些小问题，拖进去后会出现error，不影响观测，以后再解决吧。
参考知乎tensorflow profiling工具简介——nsight系列,后面还有nsight compute，下次再学着用吧。

NVTX

使用nvtx，由于我在conda environment下，所以需要安装nvtx

conda install -c conda-forge nvtx

然后在python文件中

import nvtx
nvtx.push//pop_range()#push命名，pop收集结束

最后

/usr/local/cuda-11.1/bin/nsys profile -t nvtx,osrt --force-overwrite=true --stats=true --output=quickstart python xxxx.py

即可在nsight system UI界面查看使用nvtx标记的部分
还会生成一个.sqlite文件，感觉没啥用，可以去sqlite online viewer查看具体内容。
同时也可以使用

with nvtx.annotate("name",color=""):

一般这两种方法就够了，具体可参考nvtx api。

结合launch.py

最近在想怎么profiler DGL launch.py的不同线程，想通过nsight system来查看，由于是多线程，通过上面链接通过nvtx查看server端和client端的情况，具体command

MKL_NUM_THREADS=1 OPENBLAS_NUM_THREADS=1 OMP_NUM_THREADS=1 /usr/local/cuda-11.1/bin/nsys profile -t nvtx,osrt,cuda --stats=true --force-overwrite=true --output=multiprocessing python launch.py "/home/XXXXXXX/anaconda3/envs/torch/bin/python train_dist.py --graph_name reddit --ip_config ip_config.txt --num_gpus 1 --local_rank 0 --num_epochs 1 --batch_size 1000"

Nsight Compute

CUDA C 编程/Nsight可视化相关推荐

C与CUDA混合编程的配置问题
C与CUDA混合编程的配置问题 2015-10-05 19:58 526人阅读评论(2) 收藏举报分类: CUDA(6) 目录(?)[+] 原文: http://blog.csdn.net/u0 ...
CUDA C编程接口技术分析
CUDA C编程接口技术分析编程接口 CUDA C为熟悉C编程语言的用户提供了一个简单的路径,可以方便地编写程序供设备执行. 它由C语言的最小扩展集和运行库组成. 核心语言扩展已经引入:cuda c ...
CUDA C++编程接口：编译
CUDA C++编程接口:编译一．概述 CUDA C++为熟悉C++编程语言的用户提供了一个简单的路径,以方便地编写程序以执行该设备. 它由一组最小的扩展到C++语言和运行库. 在编程模型中引入了核 ...
CUDA C++编程手册（总论）
CUDA C++编程手册(总论) CUDA C++ Programming Guide The programming guide to the CUDA model and interface. C ...
《CUDA C编程权威指南》——1.5节总结
本节书摘来自华章社区<CUDA C编程权威指南>一书中的第1章,第1.5节总结,作者[美] 马克斯·格罗斯曼(Max Grossman) ,更多章节内容可以访问云栖社区"华章社区 ...
《CUDA C编程权威指南》——3.4　避免分支分化
本节书摘来自华章计算机<CUDA C编程权威指南>一书中的第3章,第3.4节,作者 [美] 马克斯·格罗斯曼(Max Grossman),译颜成钢殷建李亮,更多章节内容可以访问云栖社 ...
c cuda 指定gpu_《CUDA C编程权威指南》——1.3　用GPU输出Hello World-阿里云开发者社区...
本节书摘来自华章计算机<CUDA C编程权威指南>一书中的第1章,第1.3节,作者 [美] 马克斯·格罗斯曼(Max Grossman),译颜成钢殷建李亮,更多章节内容可以访问云栖社 ...
Gtk与Cuda混合编程
大家都知道,不同的编译器有不同的-L和-I,而在编译过程中要实现两种编译器的混合使用,利用Makefile是一个很好地解决方案,比如之前在mpi与cuda混合编程中讨论的那样,基于图形界面的GTK编程 ...
《CUDA C编程权威指南》——2.2　给核函数计时
本节书摘来自华章计算机<CUDA C编程权威指南>一书中的第2章,第2.2节,作者 [美] 马克斯·格罗斯曼(Max Grossman),译颜成钢殷建李亮,更多章节内容可以访问云栖社 ...
【OpenCV CUDA】OpenCV和Cuda结合编程
一.利用OpenCV中提供的GPU模块目前,OpenCV中已提供了许多GPU函数,直接使用OpenCV提供的GPU模块,可以完成大部分图像处理的加速操作. 基本使用方法,请参考:http://www ...

CUDA C 编程/Nsight可视化