CUDA系统拆解-11-经典CUDA算子模式：elementwise、reduction、reorder 与 blocked compute

本文是「CUDA系统拆解」系列第 11 篇。
系列导读：CUDA系统拆解-00-导读：从编程模型到 AI 推理系统的学习路线
上一篇：CUDA系统拆解-10-Profiling、调试与瓶颈定位：先找到根因再谈优化
下一篇：CUDA系统拆解-12-面试总复盘：怎么把 CUDA 能力连接到 AI 推理

1. 这篇解决什么问题

经典 CUDA kernel 到底能分成哪几种模式。
elementwise、reduction、reorder、blocked compute 各自怎么做线程映射。
reduction、transpose、GEMM 这些典型算子，瓶颈为什么完全不同。
看一个 kernel 时，怎样从 work split、访存和同步去读它。
怎样把这些模式直接映射到 AI 推理里的常见算子。

2. 先记住的核心结论

经典 CUDA 算子大致可以分成五类：elementwise、reduction、reorder、blocked compute、composite。
elementwise 通常最像“一线程一个元素”，常见瓶颈是 memory-bound。
reduction 的关键不是公式，而是“如何把很多线程的中间结果安全、高效地合并起来”。
reorder 的重点不在计算，而在访存模式和数据重排，transpose 是最典型案例。
blocked compute 的核心是 tile 化和数据复用，GEMM 是最重要代表。
很多 AI 推理算子都不是全新物种，而是这些经典模式的组合。

3. 正文讲解

3.1 先建立一张算子模式地图

看到一个 CUDA kernel，先别急着看每一行代码，先判断它属于哪种模式。

elementwise

每个线程处理一个或几个独立元素
线程之间几乎没有依赖
常见于向量加法、激活函数、逐元素量化/反量化

reduction

多个线程共同算一个结果
必须合并中间结果
常见于 sum、max、norm、softmax 的分母计算

reorder

核心不是算，而是搬和重排
常见于 transpose、layout transform、某些 attention 前后的数据排布

blocked compute

一个 block 负责输出空间的一块 tile
输入数据通常会先进入 shared memory
常见于 GEMM、卷积、attention 的核心矩阵块计算

composite

上面几种模式的组合
常见于 softmax、LayerNorm、RMSNorm、一些 fused kernel

这一张图很重要，因为后面你看任何 kernel，都可以先把它放进某个框里，再谈瓶颈和优化。

3.2 模式一：`elementwise`

最基础的骨架长这样：

__global__ void vecAdd(const float* A, const float* B, float* C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

这个模式最重要的特征是：

问题空间通常是一维或容易摊平成一维
一个线程负责一个输出位置
线程之间几乎没有同步需求

它的性能直觉也最简单：

算得不多
读写不少
常常是典型的 memory-bound

所以 elementwise kernel 的关键通常不是复杂同步，而是：

线程映射是否自然
访存是否连续
是否值得做 fusion，减少中间结果回写

3.3 模式二：`reduction`

最小代表性骨架可以写成：

__global__ void reduceSum(const float* x, float* out, int N) {
    __shared__ float buf[256];
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + tid;

    buf[tid] = (idx < N) ? x[idx] : 0.0f;
    __syncthreads();

    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) {
            buf[tid] += buf[tid + stride];
        }
        __syncthreads();
    }

    if (tid == 0) out[blockIdx.x] = buf[0];
}

这个模式和 elementwise 最大的不同在于：

一个输出不再对应一个线程
线程之间必须协作
中间结果要分阶段合并

所以 reduction 的核心问题不再只是访存，而是：

block 内怎样组织协作
shared memory 是否合适
barrier 放得是否正确
后半程线程闲置是否严重

你要特别记住 reduction 的性能直觉：

它常常不是纯算力问题
它会同时受同步、访存、并行度下降影响

这也是为什么 reduction 很适合用来考你是否真的理解 shared memory 和 __syncthreads()。

3.4 模式三：`reorder`

transpose 是最经典代表。
它的难点不是公式，而是访存方向。

最小骨架可以这样理解：

__global__ void transposeNaive(const float* in, float* out, int W, int H) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < W && y < H) {
        out[x * H + y] = in[y * W + x];
    }
}

这个 kernel 逻辑完全正确，但经常不快。
因为它暴露了 reorder 模式的本质：

算得很少
但读写方向往往无法同时都对 global memory 友好

所以这类 kernel 的主问题通常是：

coalescing 好不好
stride 大不大
是否需要用 shared memory 做 staging 和重排

这就是为什么 transpose 常被当成“访存优化范例”，而不是“数学计算范例”。

3.5 模式四：`blocked compute`

GEMM 是这个模式里最重要的代表。

一个最小 tiled 骨架可以写成：

#define TILE 16

__global__ void gemmTiled(const float* A, const float* B, float* C,
                          int M, int N, int K) {
    __shared__ float As[TILE][TILE];
    __shared__ float Bs[TILE][TILE];

    int row = blockIdx.y * TILE + threadIdx.y;
    int col = blockIdx.x * TILE + threadIdx.x;
    float sum = 0.0f;

    for (int t = 0; t < (K + TILE - 1) / TILE; ++t) {
        int aCol = t * TILE + threadIdx.x;
        int bRow = t * TILE + threadIdx.y;

        As[threadIdx.y][threadIdx.x] =
            (row < M && aCol < K) ? A[row * K + aCol] : 0.0f;
        Bs[threadIdx.y][threadIdx.x] =
            (bRow < K && col < N) ? B[bRow * N + col] : 0.0f;

        __syncthreads();

        for (int k = 0; k < TILE; ++k) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }

        __syncthreads();
    }

    if (row < M && col < N) C[row * N + col] = sum;
}

这个模式最关键的点不是“矩阵乘法公式”，而是：

一个 block 负责输出矩阵的一块 tile
输入 tile 被 block 内线程协作搬到 shared memory
每次 global memory 读进来的数据，会被尽可能多次复用

所以 blocked compute 的核心是：

tile 切得是否合理
shared memory 是否真的提高了复用
寄存器和 shared memory 是否压得太重
block 内同步是否必要且不过多

它和前几类模式最大的区别是：

它不是在减少线程之间的关系，而是在有组织地利用线程协作和数据复用。

3.6 模式五：`composite`

这类算子通常不是单一模式，而是多个模式拼起来。

最典型的例子是 softmax：

先做 max reduction
再逐元素减 max 和 exp
再做 sum reduction
再逐元素归一化

LayerNorm / RMSNorm 也很像：

先 reduction
再 elementwise

所以 composite 模式的关键不是“它更复杂”，而是：

中间结果是否要回写 global memory
哪些阶段值得 fusion
哪些阶段适合 block 内完成
哪些地方会把 launch overhead 放大

这也是为什么很多推理 kernel 不追求“每一步都单独清晰”，而是追求“把多个阶段尽量留在同一个 kernel 里做完”。

3.7 reduction、transpose、GEMM 到底差在哪

这三个是最适合对比着学的。

reduction

重点是合并
主要矛盾是协作、同步、后半程并行度下降

transpose

重点是重排
主要矛盾是访存模式、coalescing、bank conflict

GEMM

重点是复用
主要矛盾是 tile 设计、数据流、shared memory 和寄存器的平衡

这三个模式一旦分清楚，后面很多“看起来不同”的推理算子，其实都能归到这三条主线之一。

3.8 怎样通过 `work split`、访存和同步去读一个 kernel

这是这篇最重要的实战部分。

以后你看到一个 kernel，先问四个问题：

第一，work split 是什么？

一个线程负责什么
一个 block 为什么负责这块数据
grid 怎样覆盖全部输出空间

第二，数据主要从哪来，往哪去？

是直接从 global memory 读写
还是会先经过 shared memory
中间结果是否主要留在寄存器

第三，线程之间有没有协作？

完全独立
只在 warp 内协作
还是 block 内要共享数据和同步

第四，瓶颈更像哪一类？

memory-bound
compute-bound
sync-heavy
launch-sensitive

如果你能按这四步读 kernel，很多陌生代码都会变得可拆解。

3.9 每类模式的性能直觉

elementwise

常见瓶颈：带宽
关键优化：coalescing、fusion、减少中间回写

reduction

常见瓶颈：同步、访存、后半程并行度浪费
关键优化：warp-level reduce、分层归约、减少 barrier

reorder

常见瓶颈：访存模式差
关键优化：shared memory staging、padding、layout 设计

blocked compute

常见瓶颈：数据复用不够、资源压力过高
关键优化：tiling、寄存器复用、shared memory 复用、pipeline 设计

composite

常见瓶颈：中间结果回写过多、launch 太碎
关键优化：fusion、阶段合并、就地消费中间值

4. 和 AI 推理的关系

这篇和 AI 推理的关系非常直接，因为推理里的大多数热点算子，都能拆成上面的模式。

几个典型映射：

Linear / MLP：blocked compute 为主，本质接近 GEMM
softmax、LayerNorm、RMSNorm：composite，里面带有 reduction 和 elementwise
attention：通常是 GEMM-like 块计算、reduction、reorder 的组合
KV cache 读写、layout transform：常带有 reorder 特征
激活、bias、残差：很多是 elementwise

所以你后面看推理 kernel 时，真正该做的不是把它们当成“特殊算子名词”，而是先问：

它属于哪种经典模式
它真正的瓶颈是协作、访存还是复用
它为什么会被做成 fused kernel

5. 常见误区

经典算子实战就是刷几个教学样例。不是，真正要学的是模式和判断框架。
elementwise 简单，所以没什么价值。不是，很多推理里的小算子恰恰最容易受带宽和 launch 影响。
reduction 就是“会写 shared memory 求和”。不是，重点是协作边界和合并路径设计。
transpose 只是数据搬一下。不是，它是访存重排和 shared memory 使用的经典范式。
GEMM 只是矩阵乘法。不是，它是 tile 化、复用、寄存器和 shared memory 平衡的代表。
softmax / LayerNorm 只要懂数学公式就行。不是，在 CUDA 上真正难的是阶段组织和中间结果数据流。

6. 复习自测

经典 CUDA 算子大致能分成哪几类模式？
elementwise、reduction、reorder、blocked compute 各自的线程映射方式有什么区别？
reduction、transpose、GEMM 的核心矛盾分别是什么？
看一个 kernel 时，怎样从 work split、访存和同步去读它？
为什么 elementwise 常常是 memory-bound？
为什么 transpose 常常要借助 shared memory？
为什么 GEMM 天然适合 tile 化和数据复用？
为什么很多 AI 推理算子本质上都是这些经典模式的组合？

系列导航

导读：CUDA系统拆解-00-导读：从编程模型到 AI 推理系统的学习路线
上一篇：CUDA系统拆解-10-Profiling、调试与瓶颈定位：先找到根因再谈优化
下一篇：CUDA系统拆解-12-面试总复盘：怎么把 CUDA 能力连接到 AI 推理
延伸阅读：CUDA实践-02-归约与线程粗化
延伸阅读：CUDA实践-04-矩阵乘优化与寄存器分块