零、前言：
此demo的Github仓库（持续更新）：github.com/Conqueror71…
📕欢迎访问：
Bilibili：space.bilibili.com/57089326
有任何疏忽和错误欢迎各位读者指出！
⭐有关性能分析的部分可直接跳转至第五部分~
一、CUDA是什么
CUDA是一种通用的并行计算平台和编程模型，是在C语言基础上扩展的。
CUDA程序的执行过程：
- 分配CPU内存，并进行数据初始化；
- 分配GPU内存，把数据从CPU内存拷贝到GPU内存；
- 调用核函数对存储在GPU内存中的数据进行运算；
- 将数据s从GPU内存拷贝到CPU内存；
- 释放CPU和GPU上分配的内存。
GPU需要与CPU协同工作才行，所以说GPU并行计算，事实上指的是基于CPU+GPU的异构计算架构，GPU和CPU通过PCIe总线连接在一起来协同工作，示意图如下：
我们发现，GPU有着更多的ALU（绿色部分），所以很适合数据并行的计算密集型任务，例如张量运算等。
CPU没有很多ALU，但是可以实现复杂的逻辑运算适合控制密集型任务。
二、CUDA编程模型的两大特色
1 通过层次结构组织线程
说到线程，我们先来明确一下进程、线程和协程的区别：
进程：是并发执行的程序在执行过程中分配和管理资源的基本单位，是一个动态概念，竞争计算机系统资源的基本单位。
线程：是进程的一个执行单元，是进程内科调度实体。比进程更小的独立运行的基本单位。线程也被称为轻量级进程。
协程：是一种比线程更加轻量级的存在。一个线程也可以拥有多个协程。其执行过程更类似于子例程，或者说不带返回值的函数调用。
如果是对于CPU而言，一个核大概只能支持1~2个硬件线程，但是对于GPU而言，在硬件层面上就能同时支持千百个并发线程，在GPU编程的时候需要考虑到这一点，以提高运行效率。
在CUDA编程中，线程的管理层次有如下4种：
- 线程网络Grid
- 线程块Block
- 线程束Warp
- 线程Thread
当核函数在主机端也就是CPU启动时，其执行会移动到设备GPU上，此时设备GPU中会产生大量的线程并且每个线程都执行由核函数指定的语句。
2 通过层次结构组织内存
如何组织内存和使用内存，是提高效率的关键，GPU中的各级缓存以及各种内存是可以软件控制的，在编程时我们可以手动指定变量存储的位置。
内存会有很多种，主要有：
- 寄存器
- 共享内存
- 常量内存
- 全局内存
为了发挥这些不同的内存的功效，在CUDA编程中会有很多的little tips，例如：
- 尽量使用寄存器
- 尽量将数据声明为局部变量
- 在重复利用某一数据时将其放入共享内存中
- 注意合并访问某些全局数据以减少访问需要的次数
GPU内存层次结构的示意图如下：
一些说明：
- 寄存器Registers，带宽=8TB/s，delay=1时钟周期，特点是快。
- 共享内存Shared Memory，带宽=1.5TB/s，delay=1~32时钟周期，特点是可受用户控制。
- 全局内存Global Memory，是GPU中最大、延迟最高并且最常使用的内存。
三、Kernel核函数
- 核函数用
__global__符号声明； - 核函数在调用时要用
<<<grid, block>>>来指定Kernel要执行的线程数量； - 每个线程都要执行Kernel，并且会分配到一个唯一的线程号
thread ID，其值可通过核函数的内置变量threadIdx来获得。
因为GPU是异构模型，所以在CUDA中通过函数类型限定词区分host和device上的函数：
__global__：在device上执行，从host中调用，返回类型必须是void，不支持可变参数，不能成为类成员函数。注意用__global__定义的kernel是异步的，这意味着host不会等待kernel执行完就执行下一步；__device__：在device上执行，仅可以从device中调用，不可以和__global__同时用；__host__：在host上执行，仅可以从host上调用，一般省略不写，不可以和__global__同时用，但可和__device__，此时函数会在device和host都编译。
2-dim线程的组织结构示意图如下：
// 缺省的值初始化为1，此处由于使用dim3，包含x, y, z三个维度，故有一个维度缺省
dim3 grid(3, 2);
dim3 block(5, 3);
kernel_fun<<< grid, block >>>(prams...);
一些说明：
-
同一网格Grid中的所有线程共享相同的全局内存空间；
-
线程网格和线程块从逻辑上代表了一个核函数的线程层次结构，这种组织方式可以帮助我们有效地利用资源，优化性能。
-
一个线程需要两个内置的坐标变量
blockIdx&threadIdx来唯一标识，它们都是dim3类型变量，其中blockIdx指明线程所在grid中的位置，而threaIdx指明线程所在block中的位置，如图中的Thread(1,1)满足：threadIdx.x = 1 threadIdx.y = 1 blockIdx.x = 1 blockIdx.y = 1
四、SM流式多处理器
一个Block上的线程是放在同一个流式多处理器SM上的，但是单个SM的资源有限，这导致线程块中的线程数是有限制的，现代GPUs的线程块可支持的线程数可达1024个。
SM是GPU硬件的一个核心组件，SM的核心组件又包括CUDA核心，共享内存，寄存器等，SM可以并发地执行数百个线程，并发能力就取决于SM所拥有的资源数。
当一个kernel被执行时，它的gird中的线程块被分配到SM上。
要注意的是，一个线程块只能在一个SM上被调度，但SM一般可以调度多个线程块。有可能一个kernel的各个线程块被分配多个SM，所以说grid只是逻辑层，而SM才是执行的物理层。
逻辑层和物理层的示意图：
由于SM的基本执行单元是包含32个线程的线程束，所以block大小一般要设置为32的倍数。
五、ltrace & strace简介
1 ltrace
Linux环境：
功能：跟踪进程的库函数调用。
与strace的对比：
- ltrace会显现出哪个库函数被调用，而strace则是跟踪程序的每个系统调用。
- ltrace与strace使用的技术大体相同，但ltrace在对支持fork和clone方面，不如strace。strace在收到fork和clone等系统调用后，做了相应的处理，而ltrace没有。
2 strace
Linux环境：
- strace命令是一个集诊断、调试、统计与一体的工具，可用来追踪调试程序，能够与其他命令搭配使用。
- Linux系统管理员可以在不需要源代码的情况下即可跟踪系统的调用。
- strace显示有关进程的系统调用的信息，这可以帮助确定一个程序使用的哪个函数，当然在系统出现问题时可以使用 strace定位系统调用过程中失败的原因，这是定位系统问题的很好的方法。
为什么要用strace？
- 在操作系统运维中会出现程序或系统命令运行失败，通过报错和日志无法定位问题根因。
- 如何在没有内核或程序代码的情况下查看系统调用的过程。
六、ltrace & nvprof实战demo
1 nvprof简介
nvprof是NVIDIA提供的一种性能分析工具，用于在Linux系统上分析CUDA应用程序的性能。它可以收集应用程序在GPU上执行时的各种指标，如内存带宽、计算吞吐量、指令执行时间等，以帮助开发人员对应用程序进行性能分析和优化。
使用nvprof工具可以获取以下信息：
- GPU 活动摘要：可以查看 GPU 活动的总体概述，如 GPU 利用率、内存带宽、计算吞吐量等；
- CUDA 核心指标：可以查看 CUDA 核心执行时间、指令执行次数、内存带宽和缓存效率等指标；
- 内存分配指标：可以查看内存分配和释放的次数和时间；
- API 指标：可以查看 CUDA API 调用的时间和次数；
- OpenACC 指标：可以查看 OpenACC 指令的执行时间和指令执行次数等。
2 ltrace实战
安装：
sudo apt-get install ltrace
执行：
ltrace -e cuda* ./hello
3 nvprof实战
安装：
sudo apt install nvidia-profiler
执行：
nvprof -o profile.json ./hello
注意，我们要把C++代码里面的new和delete给删掉，换成CUDA专用的一些内存管理方式，否则不会加速。
另外，我们如果想充分利用CUDA设备的计算资源，可以使用多个线程块和多个线程，从而达到并行计算。
例如：
int blockSize = 256;
int numBlocks = (N + blockSize - 1) / blockSize;
add<<<numBlocks, blockSize>>>(N, x, y);
当然，如果还有一个点不注意的话还是会寄，就像这样：
收到信号139通常表示程序因为内存访问错误而崩溃。
如果通过调用cudaMallocManaged函数动态分配了x和y数组的内存，但是在初始化数组时，又重新定义了x和y数组，这样就导致cudaMallocManaged分配的内存被泄漏了。
要解决这个问题，您需要将初始化数组的代码删除或注释掉，只使用cudaMallocManaged函数来分配和释放内存。
修改后的代码已经放入Github仓库github.com/Conqueror71…
随后可以正确运行，虽然还有Error，但是不影响结果，我们将在后续解决：
