CUDA 并行实现 #

Thread —— 最小逻辑并行单元 #

CUDA 里的 Thread 可以先理解成最小的并行单元

比如这个 kernel：

__global__ void add_one(float* x, float* y, int n) {
    int i = threadIdx.x;

    if (i < n) {
        y[i] = x[i] + 1.0f;
    }
}

如果你启动 10 个 thread，那么每个 thread 都会执行这一段代码。

区别在于：

threadIdx.x

每个 thread 的值不同。

例如：

thread 0: threadIdx.x = 0
thread 1: threadIdx.x = 1
thread 2: threadIdx.x = 2
...
thread 9: threadIdx.x = 9

所以虽然所有 thread 执行同一份代码，但它们处理的数据不同。

这就是 CUDA 的基本思想。

为什么还需要 Block #

如果只有 Thread，会有一个问题：

一次 kernel 可能要启动几百万、几千万个 thread。

CUDA 不会把所有 thread 扔成一大坨，而是把它们分组，这个组就叫 Block。

可以理解成：

Thread：单个学生
Block：一个班级
Grid：整个年级

比如我们启动：

kernel<<<3, 4>>>(...);

意思是：

3 个 Block
每个 Block 里有 4 个 Thread

所以总 thread 数是：

3 × 4 = 12 个 thread

threadIdx.x 只是在 当前 Block 内部编号。

所以 Block 0 里面有一个 threadIdx.x = 0

Block 1 里面也有一个 threadIdx.x = 0

全局编号：

int i = blockIdx.x * blockDim.x + threadIdx.x;

其中：

blockIdx.x：当前是第几个 Block
blockDim.x：每个 Block 有多少个 Thread
threadIdx.x：当前 Thread 在 Block 内部的编号

假设：

kernel<<<3, 4>>>(...);

那么：

gridDim.x = 3
blockDim.x = 4

SM 是什么 #

Streaming Multiprocessor，流式多处理器，对应逻辑上的 Block。

刚才讲的是 CUDA 的软件抽象，现在看硬件真实执行。

假设 GPU 有 2 个 SM

启动：

kernel<<<3, 4>>>(...);

也就是 3 个 Block：

Block 0
Block 1
Block 2

GPU 可能这样调度：

SM 0 执行 Block 0
SM 1 执行 Block 1
Block 2 先等着

等某个 SM 空出来后：

SM 0 执行完 Block 0
SM 0 再执行 Block 2

你不需要写：

Block 0 去 SM 0
Block 1 去 SM 1
Block 2 去 SM 0

这是硬件调度器自动做的。

这个地方就要针对硬件级别进行优化。

Warp 是什么 #

Warp 是 GPU 硬件调度和执行的最小物理单元，通常包含 32 个 Thread。

也就是说，你写的是 Thread，但 GPU 不会真的一个 thread 一个 thread 调度，它会把 thread 打包成 Warp。

比如你写：

kernel<<<1, 256>>>(...);

意思是：

1 个 Block
256 个 Thread

硬件会把这 256 个 Thread 分成：

256 / 32 = 8 个 Warp

也就是：

Warp 0: thread 0 ~ 31
Warp 1: thread 32 ~ 63
Warp 2: thread 64 ~ 95
...
Warp 7: thread 224 ~ 255

你不用手动创建 Warp。

Warp 是硬件自动形成的。

• Thread 是你写代码时看到的逻辑单位
• Warp 是硬件真正调度执行的单位

最好尽量让 CUDA 线程组织"契合 Warp" —— 32 的倍数，假如 33：

Thread 0  ~ 31  → Warp 0，满的
Thread 32       → Warp 1，只有 1 个有效线程

SP 又是什么 #

SP 是真正进行加减乘除的计算执行单元。

GPU 内部会把 Warp 的指令发给执行单元去算。

所以关系是：

你控制：Grid / Block / Thread

硬件控制：Block 放到哪个 SM，Thread 怎么组成 Warp，Warp 怎么用 SP 执行

大致映射关系是：

一个 Kernel 启动一个 Grid 
一个 Grid 包含很多 Block 
一个 Block 会被调度到某个 SM 上执行 
一个 Block 里的 Thread 会被硬件切成多个 Warp 
Warp 才是硬件真正调度执行的单位 
Warp 内的指令最终由 SM 内部执行单元完成

CUDA 软件模型：

Grid
 └── Block
      └── Thread

GPU 硬件执行模型：

GPU
 └── SM
      └── Warp
           └── 各种 Core / 执行单元

最重要的是：

Block → SM
Thread → Warp

这不是手动指定的，是硬件自动安排的。

全局坐标计算 #

一维坐标计算 #

一维数组 kernel，基本都是这个模板：

__global__ void kernel(float* x, float* y, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    // 一定要写边界，可能会出现空余的线程 （比如 3x4 > 10）
    if (i < N) {
        y[i] = ...;
    }
}

主机端启动：

int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize; // 向上取整，覆盖所有的数据

kernel<<<gridSize, blockSize>>>(x, y, N);

二维坐标计算 #

二维数据常见于图像、矩阵和二维 Tensor。

比如一张图像有：

width  = 图像宽度，也就是每一行有多少个像素
height = 图像高度，也就是一共有多少行

二维 kernel 里通常先计算当前线程负责的列和行：

// 列方向，当前块之前的数量 * 每块宽度 + 当前块内的线程偏移
int col = blockIdx.x * blockDim.x + threadIdx.x;
// 行方向，当前块之前的数量 * 每块高度 + 当前块内的线程偏移
int row = blockIdx.y * blockDim.y + threadIdx.y;

但是内存本质是一维连续空间，不是真正的二维格子。 所以还要把二维坐标 (row, col) 转成一维线性地址：

// 我们只需要利用行优先规则转换
if (row < height && col < width) {
    int global_idx = row * width + col;            // 转为一维线性地址
}

这里的关键是：

width 是每一行有多少个元素。
row * width 表示跳过前面 row 行。
+ col 表示当前行内部的列偏移。

例如：

width = 4
height = 3

二维数据可以看成：

第 0 行: (0,0) (0,1) (0,2) (0,3)
第 1 行: (1,0) (1,1) (1,2) (1,3)
第 2 行: (2,0) (2,1) (2,2) (2,3)

如果当前线程负责：

row = 2
col = 1

那么一维下标是：

global_idx = row * width + col;
           = 2 * 4 + 1;
           = 9;

在 CUDA 里，坐标计算不仅影响代码是否正确，也会影响性能。

一个很重要的黄金法则是：

让 x 维度的线程连续访问连续的内存地址。

原因是：在 CUDA 中，一个 Warp 通常包含 32 个 Thread，而这些 Thread 在线程编号上是连续的。对于二维 Block 来说，Warp 内的线程通常会优先沿着 threadIdx.x 方向连续排列。

也就是说，如果我们这样写：

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

int idx = row * width + col;

那么同一个 Warp 内相邻的 Thread 很可能访问：

thread 0  -> data[row * width + col + 0]
thread 1  -> data[row * width + col + 1]
thread 2  -> data[row * width + col + 2]
...
thread 31 -> data[row * width + col + 31]

这些地址在内存中是连续的。

这对 GPU 非常友好，因为相邻 Thread 访问相邻地址时，GPU 可以更高效地合并这些 global memory 访问，也就是所谓的 coalesced memory access.

因此，对于行优先存储的数据，比如 C/C++ 数组、图像、矩阵、PyTorch contiguous tensor，通常应该让：

x / col 方向对应内存中连续变化的维度

也就是让 threadIdx.x 负责访问连续的列方向。

例如二维图像中，推荐：

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

int idx = row * width + col;

而不是让相邻线程跨行、跨 stride 访问：

int idx = col * height + row;  // 对行优先数据通常不友好

可以把这个规则记成一句话：

CUDA 里不仅要让线程“算对自己的位置”，还要让相邻线程“访问相邻的数据”。

这就是为什么二维图像处理中常见的 Block 配置会选择：

dim3 block(16, 16);

或者：

dim3 block(32, 8);

它们的总线程数都是 256，都是 32 的倍数，能比较好地契合 Warp。

其中 32 × 8 的好处是：x 方向正好有 32 个线程，一个 Warp 更容易覆盖一整段连续的列方向数据，因此在行优先内存布局下更容易形成连续访存。

所以，全局坐标计算的最终目标不只是知道自己处理哪个元素

还要进一步做到：

让相邻线程处理相邻元素

这就是后续 CUDA kernel 优化的起点。