CUDA:从null到0 | scmmm's blog

本篇为 CUDA 入门的扫盲篇，以下 GPU 均默认为支持 CUDA 的 GPU，在正式写算子之前，我们先简单梳理下 CUDA 是个什么东西。

在机器学习中，存在大量简单且相似的运算，这天生适合有大量流处理器 GPU 来完成，随着 CUDA ，Tensor Core 和其他模块的完善，大量主流机器学习训练和推理任务都可以利用 GPU 加速。且常比 CPU 更快。而把他们搬到 GPU 去计算，其底层由 CUDA 算子实现。

CUDA 执行的周期

CUDA 编程模型包含 host 和 device 两个执行端。host 是 CPU 端，负责程序控制；device 是 GPU 端，负责并行计算。kernel 是运行在 device 上、由大量线程并行执行的函数。

一个典型的生命周期如下

CPU 准备数据
    -> 数据进入 GPU 可访问的内存
    -> CPU 发射 kernel
    -> GPU 线程执行
    -> CPU 或下一个 kernel 等待结果

kernel launch 相对于 CPU 默认是异步的。发射调用返回时，GPU 可能还没有完成计算；需要 cudaDeviceSynchronize、cudaStreamSynchronize 或 CUDA event 确认完成。

CUDA 执行的层次

CUDA 编程模型使用 Grid、Block、Thread，在硬件执行模型使用 Warp 取代 Thread。

1
2
3

grid
 └── block
      └── thread

Grid：一次 Kernel 启动产生的所有线程集合。
Block：线程协作单位，同一 Block 内线程可使用共享内存并同步。
Warp：GPU 实际调度和执行线程的基本单位，通常包含 32 个线程。

单个 Thread 是 Warp 中的执行元素，是编程模型中的最小线程单位。

CUDA 拓展的语法

CUDA C++ 仍然是 C++，重点只在几类扩展：

元素	作用
`__global__`	声明可由 host 发射的 kernel，返回类型必须是 void
`__device__`	声明只能由 device 代码调用的函数
`__shared__`	声明一个 block 内线程共同访问的片上内存
`__constant__`	声明 device 端只读、适合小型公共参数的内存
`<<<grid, block, shared_bytes, stream>>>`	指定 grid、block、动态 shared memory 和 stream
`threadIdx / blockIdx`	当前线程在 block 内、当前 block 在 grid 内的坐标
`blockDim / gridDim`	block 的尺寸、grid 的尺寸
`__syncthreads()`	同步一个 block 的所有线程
`atomicAdd` 等	安全地对可能被多个线程同时更新的地址做原子操作

dim3

可以理解为一个有三个分量的结构体

struct dim3 {
    unsigned int x;
    unsigned int y;
    unsigned int z;
};

kernel 启动的时候可以写成

1	kernel<<<gridSize, blockSize>>>(参数);

其中，gridSize, blockSize 都是 dim3 类型。

对于一个典型的cuda三层结构

1
2
3

grid
 └── block
      └── thread

在 CUDA kernel 里与线程组织直接相关的常见内置变量，有这几组比较核心的变量。

threadIdx
blockIdx
blockDim
gridDim

他们都是 dim3 类型的。

函数修饰符

写法	编译位置	调用方	说明
`__global__`	device	host 或 device	kernel 入口，必须返回 void，通过 <<<...>>> 调用
`__device__`	device	device	普通 device 函数，只能由 device 代码调用
`__host__`	host	host	普通 CPU 函数，默认就是 host
`__host__ __device__`	两边	两边	同一个函数生成 host 和 device 两份代码

__global__ 函数通常不能返回计算值，因为一次调用对应很多线程。结果应写入 device memory，通过输出指针传回。

A.2 Kernel launch 的 <<<...>>>

1	kernel<<<grid_dim, block_dim>>>(arg0, arg1);

写成完整形式：

1	kernel<<<grid_dim, block_dim, dynamic_shared_memory_bytes, stream>>>(args);

grid_dim：grid 中的 block 数量，可以是整数或 dim3。
block_dim：每个 block 的线程数，可以是整数或 dim3。
dynamic_shared_memory_bytes：每个 block 额外申请的动态 shared memory 字节数，默认 0。
stream：操作加入的 CUDA stream，默认是 nullptr 或默认 stream。

A.3 内建索引变量

这些变量只在 device 函数中有效，都是带 .x/.y/.z 的三维向量；未指定维度默认是 1，索引从 0 开始。

变量	含义	同一个 block 内是否相同
`threadIdx`	当前线程在 block 内的坐标	否
`blockDim`	当前 block 在 x/y/z 方向的线程数	是
`blockIdx`	当前 block 在 grid 内的坐标	是
`gridDim`	当前 grid 在 x/y/z 方向的 block 数	是

一维索引

该线程在全局线程编号即为

1	int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;

二维索引

此时先转化成二维的情况，然后

int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;

// 转换为一维线性索引（行主序）
int linearIdx = row * width + col;

三维索引

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
int index = (z * height + y) * width + x;

大概先学这么多吧，以后遇到什么不会的再补。