Nvidia 官方CUDA课程学习笔记

之前心血来潮学习了一下Nvidia CUDA，外行，文章有理解不当的地方，望指正。

主要根据以下Nvidia官方课程学习：
https://www.bilibili.com/video/BV1JJ4m1P7xW/?spm_id_from=333.337.search-card.all.click&vd_source=c256dbf86bb455b27cf32825685af0e6

一些参考链接：
https://developer.nvidia.com/zh-cn/blog/nvidia-hopper-architecture-in-depth/
https://zhuanlan.zhihu.com/p/597695921
https://www.cnblogs.com/shuimuqingyang/p/15846584.html
https://zhuanlan.zhihu.com/p/34587739
https://zhuanlan.zhihu.com/p/488340235?utm_oi=889446526974320640
https://blog.csdn.net/gzq0723/article/details/107343365

01 CUDA C Basics

CUDA（Compute Unified Device Architecture）是NVIDIA公司开发的一种编程模型和软件环境，它允许开发者使用C、C++、Python等高级语言来编写GPU程序。

CUDA 架构

使GPU具备并行计算的能力
为了高性能

CUDA C++

工业标准C++(2014)
提供扩展集合来使能异构计算
更直接的API来管理设备和内存等

异构计算(Heterogeneous Computing)(以下为个人理解)
在同一个业务中（代码中），代码分别运行在两个不同的体系架构中，Host 部分运行在CPU(常用的操作系统: Windows，Linux)，Device 部分运行在GPU(显卡或专用加速芯片)。这两部分各有各的资源，CPU常见的资源有内存，磁盘，IO设备等；GPU有内存等。
异构为CPU和GPU两部分运算

CPU 和 GPU之间使用传统的PCI-E接口或 NVLink 专用接口连接。
硬件连接架构
NVCC这个专用的编译器有两个体系的编译器，Host端的代码与平常的C/C++没有区别，NVCC会使用标准C/C++编译来编译，例如Linux下的GCC；有特别的关键字和语法来标识这是Device端的代码，NVCC会选用CUDA专用的编译器编译，如下图所示。

// 用__global__ 标识 Device 代码函数定义
__global__ void add(int *a, int *b, int *c) {
    c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x]
}

#define N 512
int main(void) {
    int *a, *b, *c;              // host 端
    int *d_a, *d_b, *d_c;        
    int size = N * sizeof(int);

    // 专用的API在Device端申请内存
    cudaMalloc((void **)&d_a, size);
    cudaMalloc((void **)&d_b, size);
    cudaMalloc((void **)&d_c, size);

    // Host 端申请内存，并用随机数初始化
    a = (int *)malloc(size); random_ints(a, N);
    b = (int *)malloc(size); random_ints(b, N);
    c = (int *)malloc(size);

    // 从Host 复制数据到 Device
    cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);

    // 调用Device端(GPU)中的函数，运算
    add<<<N, 1>>>(d_a, d_b, d_c);

    // 从Device端复制到Host端
    cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);

    // Host 端回收内存
    free(a);free(b);free(c);

    // Device 端回收内存
    cudaFree(d_a);cudaFree(d_b);cudaFree(d_c);
}

以上代码，由于有__global__标识，add函数定义在Device端，其它代码定义在Host端。提供了一些专用的API，例如cudaMalloc，cudaMemcpy, cudaFree，个人理解这些代码跑在Host端，但是又能处理Device端的内存，有点奇怪。
在调用Device端的add函数时，使用"<<<x, y>>>"来标识Device端使用的block和Thread个数，这个概念在下面会讲到。代码运行过程基本分为三步：

从Host端(CPU)复制数据到Device端(GPU)。
Device端执行该端代码。
结果从Device端复制回Host端。

Block 与 Thread

block与thread

为了并行运算，把资源分成多个Block, 每个Block又拥有多个Thread，每个Thread就是独立的运行单元。例如两个矩阵的加法运算，两个长度都是M，用传统方式可能是for循环里面运行了M次，时间上也是M次。如果用并行的方法，为每一个相加的运算分配一个Thread，那样时间上只是1次。

__global__ void add(int *a, int *b, int *c, int n) {
    // 有点像把二维数组的下标转为一维数据下标，这个下标就对应着线程ID
    int index = threadIdx.x + blockIdx.x + blockDim.x;
    if (index < n)
        c[index] = a[index] + b[index]
}

// 调用函数时指定了Block 个数和Thread个数
add<<<(N + M-1) / M, M>>>(d_a, d_b, d_c, N);

疑问：

代码有跑在CPU部分和GPU部分，哪些代码跑CPU，哪些跑在GPU，那么怎么分配合理？
在CPU的代码调用GPU的代码，会在那里死等，执行完了再返回？或者用户自己用多线程跑？或者回调？
有pci-e和nvlink两种接口，中间传输的时间导致的延迟怎么办？
答：视频里面肯定了这个延迟，肯定会比内存(Device 则)到GPU或内存(host 则)到CPU慢的。这个不在本节的讨论范围。
视频里面提到智能指针的问题，貌似并不完全肯定支持，特别是数据传输时。
答：视频意思是支持c++ 2014语言标准，但是标准库就不一定了。
CUDA 里面有 Thread, block, grid 的概念，这跟一个显卡有多少个CUDA核，或PC里面的进程，线程有什么区别和联系？
视频里面有人提到，有的编译器long是4字节，有的是8字节，cuda怎么处理这种问题？
答：应该在最初的配置里指定了与host端匹配的编译器，这样Device端与Host端的类型是一致的。

02 CUDA Share Memory

share_memory_radius

个人理解：
在许多应用场景，输入的数据有可能被下一次的计算重用到，例如窗口的移动，或者需要一些额外的数据参与计算，这时就需要Share Memory。它类似于X86 CPU里面的一级，二级缓存，是内嵌在GPU里面的。区别于Global Memory，它是在GPU外面，目前通常是DRAM，所以速率上有较大区别。Share Memroy的大小通常比较小，限制在48k bytes per thread block，以block为申请单位，block内线程共享，block间独立。

__global__ void stencil_1d(int *in, int *out){
    __shared__ int temp[BLOCK_SIZE + 2 * RADIUS];
    int gindex = threadIdx.x + blockIdx.x * blockDim.x;
    int lindex = threadIdx.x + RADIUS;

    // Read input elements into shared memory
    temp[lindex] = in[gindex];
    if (threadIdx.x < RADIUS) {
        temp[lindex - RADIUS] = in[gindex - RADIUS];
        temp[lindex + BLOCK_SIZE] = in[gindex + BLOCK_SIZE];
    }

    // Synchronize (ensure all the data is available)
    __syncthreads();

    // Apply the stencil
    int result = 0;
    for (int offset = -RADIUS; offset <= RADIUS; offset++) {
        result += temp[lindex + offset];
    }

    // Store the result
    out[gindex] = result;
}

要注意，block内的线程要做同步，例如上面代码的"__syncthreads();"。

一些问题

如果有两个线程同一时间读取同一位置的数据，会怎么样？
答：无论是全局内存还是共享内存，都会发生广播。性能不会有影响，功能就要程序员自己保证了。

共享内存支持静态方式和动态方式申请，性能没有区别。

03 CUDA Fundamental Optimization (CUDA 基础架构与优化)

3.1 GPU 与 CPU区别

CPU与GPU区别
我们经常看到CPU，GPU有同样的参数，往往都是GPU碾压CPU，例如晶体管的数量，核心数。是不是GPU厂商比CPU厂商强太多？答案是二者没啥可比性。
上图是一个经典的CPU与GPU对比图，能直观大概看出二者的区别，右边的GPU有众多的逻辑运算单元，控制管理单元也比CPU多，但是CPU的控制管理单元和逻辑运算单元都比较强大。GPU的核只能处理有限的几个指令，最常见是浮点运算（通常16位，32位，64位还分开不同的核），而CPU能处理的指令远比GPU复杂得多。控制单元也是，CPU是负责整个计算机系统的，指令的流式处理过程也比GPU复杂得多，例如GPU应该没有中断的处理。
最简单最核心的例子是，GPU某个核只能做32位浮点数加减乘除，另一个核只做向量运算。而CPU的一个核，从main函数开始到退出，能把你整个程序执行完。
GH100 流式多处理器
上图是H100的架构图，看上面绿色的小方块，INT32代表整形运算单元，FP32就是32位浮点数运算单元，FP64就是64位浮点运算单元，即程序里面的doule类型。而Tensor Core是最近几代产品专门为AI设计的运算单元，更适合向量和矩阵的运算。
每个张量核提供一个 4x4x4 矩阵处理数组，该数组执行运算 D = A * B + C ，其中答：， B 、 C 和 D 是 4 × 4 矩阵，如下图所示。矩阵乘法输入 A 和 B 是 FP16 矩阵，而累加矩阵 C 和 D 可以是 FP16 或 FP32 矩阵。
矩阵运算
张量核对 FP16 输入数据进行 FP32 累加运算。对于 4x4x4 矩阵乘法， FP16 乘法会产生一个全精度的结果，该结果在 FP32 运算中与给定点积中的其他乘积累加，如下图所示。
矩阵运算流量
H100的部分参数如下：
H100部分特性参数
上面提到的FP32 CUDA 内核，张量核(Tensor Core)相当于CPU里面的ALU，甚至远不如后者复杂，拿二者的核心数比较没有意义。

3.2 grid, block, warp(32 threads), 与 cuda 物理核的关系

从第一章就了解到，最小的执行单元是Thread, 多个Threads组成一个Block，多个Block再组成一个Grid，而每32个Threads就叫一个warp，而这些都是逻辑上的概念，软件和编程时需要遵循这个模式，物理上有一套硬件体系与之对应，但并不相等。
CUDA编程逻辑和物理对应关系
GPU硬件的一个核心组件是SM(Streaming Multiprocessor，流式多路处理器)。SM的核心组件包括CUDA核心，共享内存，寄存器等，SM可以并发地执行数百个线程，并发能力就取决于SM所拥有的资源数。

当一个kernel被执行时，它的Gird中的线程块被分配到SM上，一个线程块只能在一个SM上被调度，SM一般可以调度多个线程块。grid只是逻辑层，而SM才是执行的物理层。SM采用的是SIMT (Single-Instruction, Multiple-Thread，单指令多线程)架构，基本的执行单元是线程束(warps)，线程束包含32个线程，这些线程同时执行相同的指令，但是每个线程都包含自己的指令地址计数器和寄存器状态，也有独立的执行路径。所以尽管线程束中的线程同时从同一程序地址执行，但是可能具有不同的行为。
warp概念
当Block被划分到某个SM上时，它将进一步划分为多个线程束(warps)，因为这才是SM的基本执行单元，但是一个SM同时并发的线程束数是有限的。这是因为资源限制，SM要为每个线程块分配共享内存，而也要为每个线程束中的线程分配独立的寄存器。由于SM的基本执行单元是包含32个线程的线程束，所以Block大小一般要设置为32的倍数。

04 GPU 体系架构

4.1 SM(Streaming Multiprocessor，流式多路处理器)

如前面介绍，GPU架构的核心是并发高速运算，再根据业务配置不同的逻辑核，例如整形，浮点，矢量运算核，SM和内存是两大关键。拥有共享内存，调度器，寄存器等资源，负责指令执行的生命周期，线程(warp)调度，内存资源管理等功能。一个SM可同时包含多个Block，一个Block只能存在同一个SM。
GPU架构图(Nvidia Tesla)
在开发者角度，一个block就是一组thread的组合。而硬件上，每个block只能由单个SM执行（单个SM可以执行多个block，视情况而定）。SM里每32个连续的thread组成一个warp。单个warp的执行指令时是基于SIMT（Single Instruction Multi Thread）。
SM执行指令过程
在CUDA里面，指令是保序执行，例如上图，一定是I0, I1, I2这样的指令顺序。而CPU就不一定保序。
如上图，warp0(W0)花了2个clock周期执行了I0和I1，但是数据从内存复制到寄存器需要时间，还不能立即执行I2，这时SM会安排执行W1，如此类推。当数据传输到寄存器后，SM又回头执行W0。
如果能把所有指令周期都利用上，这样效率就最高，所以编程时应当使用尽量多的线程，例如512个线程以上，理想是2048个这种数量级。

4.2 内存管理

和CPU类似，GPU硬件上的存储也分为global memory、cache及register。global memory主要是储存由cudaMalloc所分配的数据，另外还包括各种constant和texture。一般在显卡销售的宣传上指的显存xxGB（比如RTX 3090 24GB，GTX 1060 6GB）就是global memory。Cache主要是储存需要经常访问的数据和指令。除此之外，共享内存（shared memory）也是存放到cache（具体是L1 cache）。而register则是储存执行中或即将执行的数据和指令。
GPU 内存架构
访问速度方面，由慢到快依序是global memory->cache->register，而容量方面则是相反。
各类存储分类比较
数据加载和输出流程如下图：

由于高度并行所带来的高数据吞吐量，大部分CUDA程序都是访存密集型的。对于访存的优化通常比并行优化还有效，甚至可以获得成倍的速度提升。访存优化方面主要涉及数据的合并访问（coalesced memory access）及利用更高速的shared memory或texture。

合并访问（coalesced memory access）
这部分的优化主要是针对global memory的访问。为了提高访存效率，NVIDIA在设计显卡时引入合并访问的设计。当程序的某个thread需要访问global memory上的某个节点时，内存控制器会将该节点周围的数据一起加载到cache上。这些在cache上多余的数据只允许被同一个warp的线程所访问。其他warp需要访问这些数据时需要内存控制器重新把数据加载到cache上。
内存没有冲突
上图无冲突对比有冲突的下图：
内存冲突