CUDA基础

cuda程序的后缀：.cu

编译：nvcc hello_world.cu

执行：./hello_world.cu

使用语言还是C++。

1. 核函数

__global__ void add(int *a, int *b, int *c) {
    *c = *a + *b;
}

核函数只能访问GPU的内存。也就是显存。CPU的存储它是碰不到的。

并且核函数不能使用变长参数、静态变量、函数指针。

核函数具有异步性。GPU无法控制CPU，CPU也不会去等GPU，所以需要同步，也就是显式调用同步函数。有些线程也是需要同步的。

编写CUDA程序：

int main(void){
    主机代码
    核函数调用
    主机代码
    return 0;
}

核函数不支持C++的iostream。

#include<stdio.h>
__global__ void hello_from_gpu(){
    printf("Hello from GPU!\n");
    __syncthreads();// 显式同步
}
int main(){
    hello_from_gpu<<<1,1>>>();// 显式调用核函数
    cudaDeviceSynchronize();// 显式同步
    return 0;
}

2. 线程块

int main() {
    int a = 1;
    int b = 2;
    int c;
    add<<<1, 1>>>(&a, &b, &c);
    return 0;
}

线程模型重要概念：

1. grid网格
2. block线程块

线程分块是逻辑上的划分，物理上线程不分块。

配置线程：<<<grid_num, block_num>>>

第一个参数代表着我们有M个线程块，第二个参数代表着我们的每个线程块中有N个线程。他们都是一维的。这昂个参数保存在内建变量(build-in variable)中。

gridDim.x: 该变量的数值等于执行配置中变量grid_num的值。

blockDim.x: 该变量的数值等于执行配置中变量block_num的值。

最大允许线程块的大小为1024。最大允许的网格大小是$2^31-1$（针对一维网格）。

实际使用中，总线程数大于实际使用的线程数能更好地利用计算资源，因为这样可以使得GPU在计算的时候内存访问同时进行，节省计算机计算的时间。使得核心一直处于计算中。

启动核函数后，CPU并不会等待核函数执行完毕，立马去执行主机中其他程序。所以我们要做的就是使得这两部分时间重叠。

3. 线程块的索引

int main() {
    int a = 1;
    int b = 2;
    int c;
    add<<<1, 1>>>(&a, &b, &c);
    return 0;
}

线程索引保存成内s建变量(build-in variable)：

1. blockIdx.x: 该变量指定一个线程在一个网格中的线程块索引值，范围0-girdDim.x-1。
2. threadIdx.x: 该变量指定一个线程在线程块中的索引值，范围0-blockDim.x-1。

线程具有唯一标识：

$$Idx = threadIdx.x + blockDim.x * blockIdx.x;$$

4. 推广到多维线程

1. CUDA可以组织三维的网格和线程块；

2. blockIdx和threadIdx是类型为uint3的变量，该类型是一个结构体，具有x,y,z三个成员（3个成员都为无符号类型的成员构成）:

3. 定义多维网格和线程块（C++构造函数语法）：

dim3 grid_num(Gx,Gy,Gz); dim3 block_num(Bx,By,Bz);

dim3 grid_num(2,2); // 等价于dim3 grid_num(2,2,1);
dim3 block_num(5,3); // 等价于dim3 block_num(5,3,1);

5. 一维网格 一维线程块

定义grid和block尺寸：

dim3 grid_num(4);
dim3 block_num(8);

调用核函数：

kernel_fun<<<grid_num, block_num>>>(...);

具体的线程索引方式如图所示。

blockIdx.x从0到3，threadIdx.x从0到7。

计算方式：

$$Idx = threadIdx.x + blockDim.x * blockIdx.x;$$

6. 二维网格 二维线程块

定义grid和block尺寸：

dim3 grid_num(2,2);
dim3 block_num(5,3);

调用核函数：

kernel_fun<<<grid_num, block_num>>>(...);

具体的线程索引方式如图所示。

blockIdx.x从0到1，threadIdx.y从0到1。

blockIdx.x从0到1，threadIdx.y从0到3。

计算方式：

$$int blockId = blockIdx.x + gridDim.x * blockIdx.y; int threadId = threadIdx.x + blockDim.x * threadIdx.y; int id = blockId * (blockDim.x * blockDim.y) + threadId;$$

7. 三维网格 三维线程块

定义grid和block尺寸：

dim3 grid_num(2,2,2);
dim3 block_num(5,3,1);

调用核函数：

kernel_fun<<<grid_num, block_num>>>(...);

具体的线程索引方式如图所示。

blockIdx.x、blockIdx.y和blcokIdx.z从0到1，

threadIdx.x、threadIdx.y从0到3，threadIdx.z从0到1。

计算方式：

$$int blockId = blockIdx.x + gridDim.x * blockIdx.y + gridDim.x * gridDim.y * blockIdx.z; int threadId= (threadIdx.z * (blockDim.x * blockDim.y) ) + (threadIdx.y * blockDim.x) + threadIdx.x; int id = blockId * (blockDim.x * blockDim.y * blockDim.z) + threadId;$$

三维网格、三维线程块如图所示：

https://github.com/user-attachments/assets/c57924c1-2157-4c73-87ea-36f6842e9eff

Reference

[1]. 权双.CUDA编程基础入门系列（持续更新）[EB/OL](2023-07-14)[2024-08-21].https://www.bilibili.com/video/BV1sM4y1x7of/?p=7&share_source=copy_web&vd_source=8b2bc57e71349607b55c9fde6b078ebd