📖 术语表 (Glossary)

按章节顺序整理;首次出现于第几章在最后列出。

GPU 硬件

中文英文含义首现
流式多处理器Streaming Multiprocessor (SM)GPU 上独立调度执行的硬件单元,每个 SM 内含若干 CUDA core / Tensor Core / 共享内存。A100 有 108 个 SM。Ch4
线程束warp32 个 thread 一组,作为 SM 的最小调度单位,同一 warp 在同一时钟周期执行同一条指令。Ch3
单指令多线程SIMTSingle Instruction, Multiple Threads — warp 内 32 个线程同步执行同一指令但操作不同数据。Ch4
线程束分歧warp divergence同一 warp 内不同线程走了 if/else 两个分支,会被串行执行,性能下降。Ch4
占用率occupancy实际活跃 warp 数 / SM 最大支持 warp 数。高 occupancy 利于隐藏内存延迟。Ch4
张量核心Tensor Core专门做矩阵乘累加 (MMA) 的硬件单元,从 Volta 开始引入,吞吐远高于 CUDA core。Ch9
CuTe 领域专用语言CuTe DSLCUTLASS 提供的 Python tile-programming DSL;与 CUTLASS CUDA C++ template API 并存,不是后者的同义词。Ch9
第五代张量核心指令TCGEN05Blackwell PTX 中用于异步 tensor-core MMA 及相关数据移动的一组指令族,常与 TMEM 配合。Ch6
计算能力Compute Capability (CC)GPU 硬件能力编号(如 sm_80 = Ampere A100, sm_90 = Hopper H100, sm_100 = Blackwell B100)。CC 决定可用的 Tensor Core 类型、shared memory 上限、warp 数等。PDF §5.1 Table 29-33 给出完整对照。Ch4
线程块簇Thread Block Cluster (CTA cluster)Hopper (CC 9.0+) 引入的新一级并行层级:多个 block 组成 cluster,共享 DSMEM 和 cluster barrier,__cluster_dims__() 声明。Ch14
分布式共享内存Distributed Shared Memory (DSMEM)cluster 内不同 block 的 shared memory 可互相 load/store,由 SM-to-SM 网络硬件支持。是 FlashAttention v3 加速的关键之一。Ch12/14
张量内存Tensor Memory (TMEM)Blackwell (CC 10.0+) 引入的独立片上存储,专为 accumulator 设计;与 register / shared memory 并列。容量 256 KB / SM。Ch14
张量内存加速器Tensor Memory Accelerator (TMA)Hopper 引入的硬件 DMA 单元,从 global memory 异步搬张量到 shared memory,单指令 1 KB-1 MB 块拷贝。代替 cp.asyncCh12/14
绿色上下文Green ContextCC 8.0+ 提供的 SM 子分区机制,把一个 GPU 切成多个独立 context,相互隔离 SM 资源,用于 multi-tenancy 推理。PDF §5.1 Table 30 给出每代最小分区粒度。Ch14
NVLinkNVLinkNVIDIA GPU 间高速互联,用于多 GPU all-reduce、P2P copy、scale-up 推理。Nsight Systems 可显示 NVLink/PCIe 通信时间线。Ch15
PCIePeripheral Component Interconnect ExpressGPU 与 CPU/root complex 之间的标准互联。没有 NVLink 或跨 root complex 时,多 GPU 通信常受 PCIe 拓扑限制。Ch15

编程模型

中文英文含义首现
核函数kernel__global__ 修饰的函数,由 host 调用、在 device 上由大量线程并行执行。Ch2
网格gridkernel 启动时的所有线程块的集合。Ch3
线程块block / CTA共享 shared memory 的一组 thread,最大 1024,分配到同一 SM。Ch3
合作组cooperative groupsCUDA 9+ 提供的灵活同步原语,可以细粒度同步 warp 内/block 内/grid 内线程。Ch7
stream异步命令队列,同一 stream 内严格顺序,不同 stream 间并行。Ch8
合作组 (warp 级)thread_block_tile<N>Cooperative Groups 提供的 N 线程子组(N=4/8/16/32),可比 warp 更细粒度同步与 shuffle。#include <cooperative_groups.h>。PDF §5.6.3.1.4。Ch7
合并组coalesced_group表示当前活跃线程的子组(divergent branch 内有用),cg::coalesced_threads()。PDF §5.6.3.1.5。Ch7
多屏障mbarriershared memory 中的硬件加速 split-arrive/wait barrier,CC 8.0+ 支持。是 memcpy_async + pipeline 的底层。PDF §5.6.1。Ch8/12
程序依赖启动Programmatic Dependent Launch (PDL)CUDA 11.8+ 让 kernel B 在 kernel A 未结束时提前开始 prologue(B 内调 cudaGridDependencySynchronize() 后阻塞)。LLM decode loop 用它隐藏 launch overhead。Ch14
CUDA 图CUDA Graph把一段固定 CUDA work 捕获成可重复 replay 的图,减少 CPU launch overhead。LLM decode 常按 batch/shape 分桶复用 graph。Ch15
点对点访问Peer-to-Peer Access (P2P)一个 GPU 直接访问另一个 GPU 显存的能力。CUDA 里用 cudaDeviceCanAccessPeercudaMemcpyPeerAsync 检查与传输。Ch15

内存

中文英文含义首现
全局内存global memoryGPU 显存,所有线程可访问,延迟最高(几百周期),容量最大。Ch5
共享内存shared memorySM 上的片上 SRAM,block 内所有线程共享,延迟接近 L1。Ch5
寄存器register线程私有的最快存储,每个 SM 有 64K~256K 寄存器。Ch5
合并访问memory coalescing一个 warp 的 32 个线程访问连续 128B 内存,硬件合并为一次内存事务。Ch5
存储体冲突bank conflict同一 warp 多线程访问 shared memory 的同一 bank,会被串行化。Ch6
流序分配器Stream-ordered Allocator (cudaMallocAsync)CUDA 11.2+ 的非阻塞分配器:alloc/free 入 stream 队列,同 stream 内严格顺序,跨 stream 异步。LLM KV pool 的标准做法。PDF §5.2.3 / §5.3.6.4。Ch13
显存池CUDA Memory Pool复用 GPU allocation 的内存池。推理服务通常为 KV cache、logits、attention workspace 预分配 pool,避免 hot path cudaMalloc/cudaFreeCh15
常驻 L2 缓存Persisting L2 CacheCC 8.0+ 允许把 L2 cache 一部分"set-aside"给热数据驻留,CUDA_DEVICE_DEFAULT_PERSISTING_L2_CACHE_PERCENTAGE_LIMIT 控制比例。Llama 权重 prefetch 用。PDF §5.2.3.7。Ch5/14
异步内存拷贝memcpy_async / cp.asyncCC 8.0+ 硬件加速 global→shared 异步拷贝,与计算 overlap。是 FlashAttention 的关键加速。PDF §5.6.2 / §5.4.4 Pipelines。Ch8/12
页锁定内存Pinned Host Memory不能被 OS 换出的 host buffer,CUDA 可用它做真正异步 H2D/D2H 传输。服务端推理常用 pinned ring buffer 做 request staging、token/logits 回传和 offload pipeline。Ch15
可分页主机内存Pageable Host Memory普通 CPU 内存,CUDA 拷贝时可能需要临时 staging 并阻塞 host 线程。适合初始化,不适合 decode hot path 的频繁小拷贝。Ch15
拷贝引擎Copy EngineGPU 上负责 H2D、D2H、P2P 等 DMA 传输的硬件路径。Nsight Systems 的 Memcpy/NVLink/PCIe row 可用来判断它是否与 compute overlap。Ch15
阻塞回读Blocking Readbackhost 用同步 cudaMemcpycudaDeviceSynchronize 等待 device 标量结果,decode loop 中会造成 GPU 空洞。Ch15

LLM 推理

中文英文含义首现
通用矩阵乘GEMMGeneral Matrix Multiplication, C = αAB + βC,LLM 80% 的计算来自 GEMM。Ch9
缩放点积注意力Scaled Dot-Product AttentionAttention(Q,K,V) = softmax(QKᵀ/√d) V,Transformer 核心算子。Ch11
在线 softmaxonline softmax分块计算 softmax 时维护当前最大值与归一化项的算法,是 FlashAttention 的基石。Ch10
键值缓存KV Cache推理时缓存历史 token 的 K/V 投影,避免每步重算。Ch13
分页 KV 缓存Paged KV Cache把请求的逻辑 token 位置通过 block table 映射到非连续物理 KV pages,以便复用、扩容并控制碎片。Ch14
NVFP4NVIDIA FP4Blackwell 支持的 4-bit block-scaled 浮点格式;讨论峰值时必须说明 scale 布局和 dense/sparse 口径。Ch9
MXFP4OCP Microscaling FP4采用共享指数/scale 的 OCP microscaling 4-bit 格式;与 NVFP4 不应混为同一编码。Ch9
有效吞吐Goodput满足 SLO(如 TTFT、TPOT、正确性或质量门槛)的请求/token 吞吐;不能只用原始 tokens/s 替代。Ch15
预填充/解码分离Prefill/Decode Disaggregation把 compute-heavy prefill 与 memory/latency-sensitive decode 放到可独立扩缩的 worker pool,并传输 KV 状态。Ch14
KV cache 量化KV Cache Quantization把历史 K/V cache 用 FP8、INT8、INT4 等低位宽保存,以降低长上下文 decode 的显存占用和 HBM 读取。必须同时评估 scale/dequant 成本和质量回归。Ch15
旋转位置编码Rotary Position Embedding (RoPE)用复数旋转方式注入位置信息,被 Llama 等主流模型采用。Ch13
分页注意力PagedAttentionvLLM 提出的 KV cache 分页管理,让显存利用率接近 OS 分页。Ch14
块表Block TablePagedAttention / paged KV cache 中从逻辑页映射到物理 KV block 的表。decode 扫描时 block table 读取会进入 L2/DRAM profile。Ch15
内部碎片Internal Fragmentation分配粒度大于实际使用量时产生的浪费。paged KV 中 page size 太大可能让最后一页或短请求浪费 KV cache 空间。Ch15
前缀缓存Prefix Cache复用相同系统提示词、模板或共享上下文已经生成的 KV cache。高质量实现应复用 paged block table,避免命中后复制 KV payload。Ch15
多查询注意力Multi-Query Attention (MQA)所有 query head 共享一组 K/V head,显著降低 decode KV cache 容量和读取字节,但模型质量和 kernel fanout 需要评估。Ch15
分组查询注意力Grouped-Query Attention (GQA)多个 query head 共享一个 KV head,是 MHA 和 MQA 的折中。长上下文 decode 中常用于降低 KV cache 带宽压力。Ch15
融合乘加Fused Multiply-Add (FMA)d = a*b + c 一条指令完成,IEEE 754-2008 要求;CUDA fmaf() 调到硬件 FFMA。Tensor Core 的 mma.sync 本质是大规模并行 FMA。PDF §5.5.1.5。Ch9
算子融合Kernel Fusion把多个相邻小 kernel 合并成一个 kernel,以减少 launch overhead 和中间 tensor 的 global memory 读写。LLM decode 常用于 residual、norm、RoPE、quant、sampling tail。Ch15
融合尾部Fused Epilogue把矩阵乘或 attention 后面的 bias、residual、activation、norm、quant 等操作合到同一个尾部路径里,减少额外 kernel 和 memory traffic。Ch15
warp 矩阵WMMA (Warp Matrix Multiply-Accumulate)nvcuda::wmma namespace 提供的 16×16×16(fp16)等固定形状的 warp 级矩阵乘,编译到 Tensor Core mma 指令。CC 7.0+。PDF §5.4.11。Ch9
warpgroup 矩阵WGMMA (Warpgroup MMA)Hopper 引入的 4-warp(128 线程)共同执行的更大矩阵乘,形状如 64×N×16;FlashAttention v3 / FlashMLA 用它。PTX 直接写。Ch14
持久核Persistent Kernel一个 kernel 跑到结束,循环处理多 work-item(不退出)。LLM 用法:decode loop 整个塞进一个 kernel,省 launch overhead,跟 CUDA Graph 互补。Ch14
Stream-K 调度Stream-K SchedulingCUTLASS 引入的 GEMM 调度:把 K 维切分布到 block 上,再 reduce;对极端形状(如 LLM decode 的 M=1)比经典 split-K 更好。Ch9/14
权重量化Weight Quantization把模型权重从 fp16/fp32 压到 int8/int4/fp8 等低位宽,decode 阶段主要收益是减少 HBM 读权重字节。Ch15
分组量化Group-wise Quantization每 group 共享 scale/zero-point 的量化方式。group 越小精度通常越好,但 scale 读取和元数据越多。Ch15
W8A16 / W4A168-bit / 4-bit Weight, 16-bit Activation权重为 8-bit 或 4-bit、激活保持 16-bit 的推理配置。LLM decode 中常用于降低权重带宽压力。Ch15
连续批处理Continuous Batching在线推理服务把不同请求动态合并到同一 decode batch 中,旧请求完成后新请求立刻补位,以提高 GPU 利用率。Ch15
分块 prefillChunked Prefill把长 prompt prefill 切成多个 GPU work chunk,在 chunk 边界插入 decode 或其他高优先级工作,以平衡 TTFT、TPOT 和 p99。Ch15
队头阻塞Head-of-Line Blocking一个长任务排在前面导致后续短任务等待。LLM 服务中常见于长 prefill kernel 挡住 decode token step。Ch15
流优先级Stream PriorityCUDA stream 的调度优先级提示。它可影响后续 work 的调度顺序,但通常不能抢占已经运行的长 kernel。Ch15
不规则批Ragged Batchbatch 内请求长度、活跃状态或 cache 页数不同。直接按最大长度 padding 会浪费 KV scan 带宽。Ch15
活跃请求压缩Active Request Compactionscheduler 维护当前仍需 decode 的 request id 列表,让 kernel 只处理活跃 slot,减少 inactive slot 与 padding 工作。Ch15
混合专家Mixture of Experts (MoE)每个 token 只路由到少数 expert 子网络计算的模型结构。推理瓶颈除了 expert GEMM,还包括 routing、dispatch/gather、load balance 和跨 GPU all-to-all。Ch15
专家路由Expert Routing根据 gate 结果把 token 分发到 expert,再把 expert 输出聚合回原 token 顺序。通常包含 prefix-sum、scatter、grouped GEMM、gather。Ch15
容量因子Capacity FactorMoE 为每个 expert 预留的 token 容量相对平均负载的倍率。过高浪费 padding 计算,过低可能 drop token 或触发重路由。Ch15
张量并行Tensor Parallelism (TP)把单层矩阵或 attention/MLP 权重切到多张 GPU 上并行计算,再用 all-reduce、all-gather 或 reduce-scatter 合并结果。推理中常受 NVLink/PCIe/NCCL overlap 影响。Ch15
全规约All-Reduce多 GPU 通信模式:每个 rank 输入一段张量,规约后每个 rank 都得到完整结果。Tensor parallel MLP/attention 常见瓶颈。Ch15
规约散射Reduce-Scatter先规约再把结果切分给各 rank,可减少每个 rank 持有的输出字节,常用于优化 tensor parallel 通信和 overlap。Ch15
推测解码Speculative Decoding用小 draft model 先提出多个 token,再由大 target model 批量验证。收益取决于 accept rate、verify batch、draft 成本和质量回归。Ch15
接受率Accept Ratespeculative decoding 中 draft token 被 target model 接受的比例。接受率低会增加 rejected-token verify waste,降低端到端收益。Ch15
Logits 处理器Logits Processor采样前对 logits 做 temperature、repetition penalty、bad-word mask、min-length mask 等规则处理。低 batch decode 中常表现为多个短 kernel。Ch15
核采样Nucleus Sampling / top-p按概率从累计质量达到 p 的候选 token 集合中采样。实现通常需要 top-k/top-p 选择、softmax 统计和随机数,容易成为 decode tail 热点。Ch15

工具链

中文英文含义首现
编译器nvccNVIDIA CUDA 编译器,把 .cu 编译为 host C++ + device PTX/SASS。Ch1
中间表示PTXParallel Thread eXecution,CUDA 的虚拟 ISA,类似 LLVM IR。Ch8
架构机器码SASS实际硬件执行的机器码,与 SM 架构强绑定。Ch8
性能分析器Nsight ComputeNVIDIA 的 kernel 级性能分析工具,输出 roofline、stall reason 等。Ch8
时间线分析器Nsight Systems系统级 timeline 分析,看 H2D/D2H/kernel/stream overlap。Ch8
GPU 性能分析GPU Profiling用 Nsight Systems / Nsight Compute / NVTX / CUPTI 等工具定位 GPU 程序瓶颈,核心问题是先证明受限资源,再改代码。Ch15
区间性能分析Range ProfilingCUPTI 以用户或自动 range 为单位采集 GPU metrics 的现行 target API;旧 Profiling API 自 CUDA 13.0 起 deprecated。Ch15
屋顶线模型Roofline Model用算术强度 (FLOP/Byte) 判断 kernel 受限于计算峰值还是内存带宽,是 GEMM、attention、decode GEMV 调优的第一性工具。Ch8/15
首 token 延迟Time To First Token (TTFT)从请求进入服务到产生第一个 token 的延迟,主要受 prefill、排队、调度和网络影响。Ch15
单输出 token 延迟Time Per Output Token (TPOT)decode 阶段每生成一个 token 的平均耗时,直接决定流式输出速度和并发吞吐。Ch15
启动空洞launch gapNsight Systems timeline 上两个 GPU kernel 之间由 CPU 调度、同步或 launch 开销造成的空闲间隔。decode 小 kernel 串最常见。Ch15
设备二进制Cubin编译后针对具体 sm_XY 的二进制 (cubin = CUDA binary),由 nvcc 嵌入 fatbin。PDF §5.2.2 JIT 一节讨论 cubin vs PTX 的加载策略。Ch8
胖二进制Fatbin (FATBIN)包含多份 cubin(不同 sm_XY)+ 一份 PTX 的归档,运行时根据 GPU 选 cubin 或 JIT PTX。nvcc -gencode arch=compute_80,code=sm_80 ... 控制内容。Ch8
即时编译JIT Compilation驱动在程序运行时把 PTX 编译成目标 GPU 的 SASS,缓存到 ~/.nv/ComputeCache。CC 升级时保证向前兼容。PDF §5.2.2.1-5.2.2.6。Ch8/14
NVTXNVIDIA Tools ExtensionnvtxRangePush/Pop 给 timeline 加自定义标签,Nsight Systems / Compute 都能识别。LLM 推理一般标 prefill / decode / sample 等阶段。Ch8/14
架构专属编译目标Architecture-Specific Target (compute_XYa)a 后缀的编译目标(如 compute_90a)允许用该架构独占的指令(WGMMA、TMA),但 cubin 不能跨 minor version 运行。PDF §5.1.2。Ch14