Evan's blog

Avatar

[CUDA 入门] 基础篇

8 min read

基础概念

Host & Device

CUDA 编程中，函数可以分为 Host Function, Device Function, Kernel Funciton 几种：

  • Host Function: 只能在 CPU 上调用和运行，用 __host__ 修饰
  • Device Function: 只能在 GPU 上调用和运行，用 __device__ 修饰，被核函数调用时会继承线程索引
  • Kernel Function: 也称为核函数，在 CPU 上调用，在 GPU 上运行，用 __global__ 修饰

Tip

函数也可以同时被 __host__ 和 __device__ 修饰，同时可以在 CPU 和 GPU 上调用和运行

__host__ void host_func() {}
__device__ void device_func() {}
__global__ void global_func() {}

线程与线程块

CUDA 编程的核心在于线程间的并行与合作，因此需要从逻辑上（同时也是物理上）将线程分组并索引。

  • Grid: 所有线程的集合
  • Block: 一组线程，有时也称为一个 CTA (Cooperative Thread Array)，一个 Block 内的线程可以快速地通过共享内存 (Shared Memory) 进行数据的交换
  • Thread: 一个线程，最小的执行单位

除此之外，还有 Warp（线程束）的概念：

  • Warp: 每 32 个线程组成一个 Warp，是最小的调度单元。每个时钟周期内，每个 Warp Scheduler 最多选取一个就绪 (Eligible) 的 Warp 发射指令，实际执行指令时总是以一个 Warp 为单位的，换句话说，一个 Warp 内的所有线程总是同步的

dim3

常用于索引线程或线程块，有 x, y, z 三个成员变量。

索引

如何对线程进行索引呢？在核函数中，可以通过 blockIdx 和 threadIdx 来进行索引：

  • blockIdx: 一个 dim3 变量，表示当前 Block 在 Grid 中的索引
  • threadIdx: 一个 dim3 变量，表示当前 Thread 在当前 Block 中的索引
  • 各个线程沿 x, y, z 方向依次连续，比如对于一个形状为 (4, 32) 的 Block, blockIdx.x = 0 的 32 个线程为 0-31 号线程，blockIdx.x = 1 的 32 个线程为 32-63 号线程，以此类推

启动核函数

在调用核函数时，需指定该核函数逻辑上的分块参数（即指定 gridDim 和 blockDim）。 CUDA 的语法为 kernel_name<<<gridDim, blockDim>>>(args)，比如：

dim3 grid_dim(4, 8, 16);
dim3 block_dim(16, 16);
kernel<<<grid_dim, block_dim>>>();

会启动一个含有 4*8*16 个 Block，每个 Block 含有 16*16 个 Thread 的核函数。

Tip

设置多维而不是一维的线程块大小主要是为了简化索引的计算，比如对于二维矩阵的操作就可以启动一个二维的线程块，而不用启动一个一维线程块后再在核函数内计算索引

举例来讲，假设要对一个二维矩阵中的每个元素加 1，可以这样写：

int mat[M][N];
__global__ void mat_incr() {
	int tidx = threadIdx.x, tidy = threadIdx.y;
	int blkx = blockIdx.x, blky = blockIdx.y;
	int x = blkx + tidx;
	int y = blky + tidy;
	mat[x][y]++;
}

Attention

启动核函数时每个 Block 最多启动 1024 个线程，这是硬件决定的 可以认为 Block 数目几乎没有限制

可以认为 GPU 会顺序地将 Block 启动，当有 SM 空缺时就将剩余未启动的 Block 分配给这个 SM.

GPU 架构

理解 GPU 的架构对编写高性能的 kernel 非常有帮助。

内存层级

Global Memory → L2 Cache → Shared Memory/L1 Cache → Register

  • 内存大小依次下降，读取速度依次上升
  • Shared Memory 和 L1 Cache 共享同一个内存条，用户分配的共享内存之外作为 L1 Cache
  • 共享内存 Shared Memory: 一个 Block 内的线程均可共享。因此，如果一个 Block 内的线程需要复用某些数据，一般会将这些要复用的数据放在共享内存中，减少对全局内存 (Global Memory) 的访问。共享内存的访问较为复杂，或许后续会单独写一篇 Blog……
  • L2 cache: 由所有 SM 共享，当访问数据量极大时，为了提高 L2 Cache 命中率，需要对每个 Block 的任务做 Block Level 的 Swizzle 映射

处理器层级

GPU → GPC → TPC → SM → CUDA Core/Tensor Core

  • GPU 含有多个 GPC，GPC 含有多个 TPC，……

编程视角层级

GPU → Cluster → Block/CTA → WarpGroup → Warp → Thread

  • Cluster 和 Warp Group 均是从 Hopper 架构才开始引入的概念，目前我没太研究清楚……

Tips

写 Kernel 之前

  • 首先需要对问题进行合理分块，将任务分解到 Block，尽量保证不同 Block 间的任务是独立的
  • 往往需要考虑或设计数据排布，保证每个 Thread 读取的数据是连续的。
    • 实际上，是保证每个 Warp 内的线程读取的数据是连续的，因为指令的发射总是以 Warp 为单位
    • 举例来讲，即使一个 Warp 内的线程因为某种 Swizzle 模式导致连续的线程读取的不是连续的数据，但是只要整个 Warp 读取的数据是连续的，就能实现合并访存

写 Kernel 时

  • 能否使用向量化访存？如 float4, float2 等，能减少发射的指令数，更好打满带宽

常用函数

Device 端

  • __syncthreads(): 保证一个 Block 内的线程在此处同步
    • 会破坏指令的流水线，因此只在必要时做同步，比如向共享内存写入数据后（且其他 Warp 的线程可能会访问这些数据）
  • 各种 shfl 函数，用于 Warp 内线程的快速通信，无需经过共享内存缓冲，直接在寄存器间操作，速度更快
    • 基本格式：T __shfl_sync(unsigned mask, T var, int srcLane)
    • __shfl_sync: 广播，所有线程从指定线程中取数据
    • __shfl_up_sync: 当前线程从编号比自己小指定值的线程中取数据
    • __shfl_down_sync: 当前线程从编号比自己大指定值的线程中取数据
    • __shfl_xor_sync: 从线程编号异或指定值的线程中取数据

Host 端

  • cudaMalloc: 在 GPU 上分配内存
  • cudaMemcpy: 将数据在 CPU 和 GPU 间拷贝

Graph View