CUDA 术语表

执行模型 · Execution Model

kernel

用 __global__ 标记、在 GPU(device)上执行、可从 CPU(host)调用、返回 void 的函数;启动一次会让很多线程并行各跑一遍它。

__global__

把函数声明为 kernel 的执行空间限定符 (execution space specifier):device 执行、host 可调用、返回类型必须 void,且不能是类成员函数。

host / device

host = CPU 及其内存;device = GPU 及其内存。两者内存空间相互独立。

SIMT

单指令多线程 (Single Instruction, Multiple Threads):一个 warp 内 32 个线程同时执行同一条指令;走不同分支的线程被屏蔽 → warp divergence。

nvcc

编译 .cu 源文件的 CUDA C++ 编译器。本课程目标硬件用 nvcc -arch=sm_80(A100)。

线程层级 · Thread Hierarchy

grid

一次 kernel 启动创建的所有 block 的集合;同一 grid 内所有 block 大小与维度相同,可 1D/2D/3D。每次启动恰好一个 grid。

block

thread block · 一组线程(A100 上最多 1024 个),整体在单个 SM 上执行不可迁移,内部可共享 shared memory 并用 __syncthreads() 同步。block 之间必须能独立执行。

thread

SIMT 模型中的基本并行单位;每个线程跑同一段 kernel 代码,靠内建索引变量定位自己负责的数据。

warp

block 内每 32 个线程组成的一组,是 SIMT 硬件调度/执行单位,lane 编号 0–31。block 线程数取 32 的倍数可避免空 lane 浪费。

SM

流多处理器 (Streaming Multiprocessor) · 执行 block 的硬件单元;一个 block 整体在一个 SM 上跑,一个 SM 可并发跑多个 block。A100 有 108 个 SM。

内建变量 · Built-in Variables

threadIdx

线程在其 block 内的索引(.x/.y/.z)。逐线程不同。

blockIdx

block 在 grid 内的索引(.x/.y/.z)。逐 block 不同。

blockDim

一个 block 内的线程数(.x/.y/.z)。对所有线程相同,启动时由执行配置决定。

gridDim

grid 内的 block 数(.x/.y/.z)。对所有线程相同,启动时确定。

warpSize

值为 32 的内建变量。注意:这是架构定义、运行时可读,而非语言保证的常量。

关键公式与启动 · Indexing & Launch

全局线程索引

跨整个 grid 唯一的线程编号。沿 x 轴:

<<< >>>

执行配置 (execution configuration) · kernel<<<numBlocks, threadsPerBlock>>>(args)。第 1 参 = block 数(grid 大小),第 2 参 = 每 block 线程数;可选第 3 参 = 动态 shared memory 字节数,第 4 参 = stream。

grid-stride loop

让固定大小 grid 处理比总线程数更大的数组:每次迭代按 blockDim.x * gridDim.x(grid 中线程总数)步进。

cudaDeviceSynchronize()

阻塞 host 直到 GPU 完成。因 kernel 启动异步而必需,也用于刷新 device 端 printf 输出。

竞态

race condition · 多个并行线程读写同一内存位置导致的 bug;如让多线程都执行"处理整个数组"的逻辑就会发生。

抽象层级 · Abstraction Layers

Triton

OpenAI 开源的 GPU kernel 语言与编译器。Python DSL（@triton.jit）以 block 级（program）视角写 kernel,编译器自动生成线程调度、shared memory 分配、访存合并。编译链为 Triton IR → LLVM IR → PTX → SASS。torch.compile 的默认后端 Inductor 输出 Triton kernel。

program

Triton 的并行执行单元,类似 CUDA 的 block。用 tl.program_id(axis) 获取编号,用 tl.arange(0, BLOCK_SIZE) 表示内部的向量化工作单元。程序员只管 block 级 tiling 策略,线程调度由编译器接管。

Liger-Kernel

LinkedIn 开源的 LLM 训练用 Triton kernel 库。预写 RMSNorm、RoPE、SwiGLU、CrossEntropy、FusedLinearCrossEntropy 等高频算子的优化 Triton kernel。一行 apply_liger_kernel_to_llama() 即可 monkey-patch HuggingFace 模型。技法:kernel fusion（合并 kernel 减少 HBM 往返）、chunking（大 tensor 切块逐算）、in-place（原地覆盖省分配）、recomputation（用计算换显存）。

kernel fusion

内核融合 · 把多个 kernel 合并成一个,中间结果不写回 HBM 直接在寄存器/shared memory 中消费。liger-kernel 的 FusedLinearCrossEntropy 和 FlashAttention 的 tiling 都是此技术的典型应用。

autotuning

自动调优 · Triton 编译器自动搜索最优的 block size、num_warps、num_stages 等参数组合。程序员用 @triton.autotune 标记搜索空间,编译器在运行时挑选快的那组。

torch.compile / Inductor

PyTorch 2.0 的默认 JIT 编译器。三段流水线:Dynamo（捕图 FX Graph）→ Inductor（图优化+降低为 loop-level IR）→ 默认输出 Triton kernel（也可选 C++/OpenXLA 等 backend）。详见 AI 编译器系列。

硬件与编译 · Hardware & Compilation

PTX

Parallel Thread Execution · NVIDIA GPU 的中间汇编语言。CUDA C++ / Triton 编译后会生成 PTX,再被 GPU 驱动 JIT 编译为 SASS（实际机器码）。

HBM

High Bandwidth Memory · GPU 的全局显存（~1.5-2 TB/s on A100）。每次读写都有数百个周期的延迟,是主要瓶颈。kernel fusion 的目的就是少读少写 HBM。

coalescing

合并访存 · 让一个 warp 的 32 个线程访问连续的全局内存地址,GPU 硬件可合并为一次大事务读写。Triton 编译器自动保证 coalescing,CUDA 里要手动保证。