| 中文 | 英文 | 含义 | 首现 |
| 流式多处理器 | Streaming Multiprocessor (SM) | GPU 上独立调度执行的硬件单元,每个 SM 内含若干 CUDA core / Tensor Core / 共享内存。A100 有 108 个 SM。 | Ch4 |
| 线程束 | warp | 32 个 thread 一组,作为 SM 的最小调度单位,同一 warp 在同一时钟周期执行同一条指令。 | Ch3 |
| 单指令多线程 | SIMT | Single 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 DSL | CUTLASS 提供的 Python tile-programming DSL;与 CUTLASS CUDA C++ template API 并存,不是后者的同义词。 | Ch9 |
| 第五代张量核心指令 | TCGEN05 | Blackwell 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.async。 | Ch12/14 |
| 绿色上下文 | Green Context | CC 8.0+ 提供的 SM 子分区机制,把一个 GPU 切成多个独立 context,相互隔离 SM 资源,用于 multi-tenancy 推理。PDF §5.1 Table 30 给出每代最小分区粒度。 | Ch14 |
| NVLink | NVLink | NVIDIA GPU 间高速互联,用于多 GPU all-reduce、P2P copy、scale-up 推理。Nsight Systems 可显示 NVLink/PCIe 通信时间线。 | Ch15 |
| PCIe | Peripheral Component Interconnect Express | GPU 与 CPU/root complex 之间的标准互联。没有 NVLink 或跨 root complex 时,多 GPU 通信常受 PCIe 拓扑限制。 | Ch15 |
| 中文 | 英文 | 含义 | 首现 |
| 核函数 | kernel | 用 __global__ 修饰的函数,由 host 调用、在 device 上由大量线程并行执行。 | Ch2 |
| 网格 | grid | kernel 启动时的所有线程块的集合。 | Ch3 |
| 线程块 | block / CTA | 共享 shared memory 的一组 thread,最大 1024,分配到同一 SM。 | Ch3 |
| 合作组 | cooperative groups | CUDA 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 |
| 多屏障 | mbarrier | shared 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 里用 cudaDeviceCanAccessPeer 和 cudaMemcpyPeerAsync 检查与传输。 | Ch15 |
| 中文 | 英文 | 含义 | 首现 |
| 全局内存 | global memory | GPU 显存,所有线程可访问,延迟最高(几百周期),容量最大。 | Ch5 |
| 共享内存 | shared memory | SM 上的片上 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/cudaFree。 | Ch15 |
| 常驻 L2 缓存 | Persisting L2 Cache | CC 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.async | CC 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 Engine | GPU 上负责 H2D、D2H、P2P 等 DMA 传输的硬件路径。Nsight Systems 的 Memcpy/NVLink/PCIe row 可用来判断它是否与 compute overlap。 | Ch15 |
| 阻塞回读 | Blocking Readback | host 用同步 cudaMemcpy 或 cudaDeviceSynchronize 等待 device 标量结果,decode loop 中会造成 GPU 空洞。 | Ch15 |
| 中文 | 英文 | 含义 | 首现 |
| 通用矩阵乘 | GEMM | General Matrix Multiplication, C = αAB + βC,LLM 80% 的计算来自 GEMM。 | Ch9 |
| 缩放点积注意力 | Scaled Dot-Product Attention | Attention(Q,K,V) = softmax(QKᵀ/√d) V,Transformer 核心算子。 | Ch11 |
| 在线 softmax | online softmax | 分块计算 softmax 时维护当前最大值与归一化项的算法,是 FlashAttention 的基石。 | Ch10 |
| 键值缓存 | KV Cache | 推理时缓存历史 token 的 K/V 投影,避免每步重算。 | Ch13 |
| 分页 KV 缓存 | Paged KV Cache | 把请求的逻辑 token 位置通过 block table 映射到非连续物理 KV pages,以便复用、扩容并控制碎片。 | Ch14 |
| NVFP4 | NVIDIA FP4 | Blackwell 支持的 4-bit block-scaled 浮点格式;讨论峰值时必须说明 scale 布局和 dense/sparse 口径。 | Ch9 |
| MXFP4 | OCP 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 |
| 分页注意力 | PagedAttention | vLLM 提出的 KV cache 分页管理,让显存利用率接近 OS 分页。 | Ch14 |
| 块表 | Block Table | PagedAttention / 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 Scheduling | CUTLASS 引入的 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 / W4A16 | 8-bit / 4-bit Weight, 16-bit Activation | 权重为 8-bit 或 4-bit、激活保持 16-bit 的推理配置。LLM decode 中常用于降低权重带宽压力。 | Ch15 |
| 连续批处理 | Continuous Batching | 在线推理服务把不同请求动态合并到同一 decode batch 中,旧请求完成后新请求立刻补位,以提高 GPU 利用率。 | Ch15 |
| 分块 prefill | Chunked Prefill | 把长 prompt prefill 切成多个 GPU work chunk,在 chunk 边界插入 decode 或其他高优先级工作,以平衡 TTFT、TPOT 和 p99。 | Ch15 |
| 队头阻塞 | Head-of-Line Blocking | 一个长任务排在前面导致后续短任务等待。LLM 服务中常见于长 prefill kernel 挡住 decode token step。 | Ch15 |
| 流优先级 | Stream Priority | CUDA stream 的调度优先级提示。它可影响后续 work 的调度顺序,但通常不能抢占已经运行的长 kernel。 | Ch15 |
| 不规则批 | Ragged Batch | batch 内请求长度、活跃状态或 cache 页数不同。直接按最大长度 padding 会浪费 KV scan 带宽。 | Ch15 |
| 活跃请求压缩 | Active Request Compaction | scheduler 维护当前仍需 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 Factor | MoE 为每个 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 Rate | speculative 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 |
| 中文 | 英文 | 含义 | 首现 |
| 编译器 | nvcc | NVIDIA CUDA 编译器,把 .cu 编译为 host C++ + device PTX/SASS。 | Ch1 |
| 中间表示 | PTX | Parallel Thread eXecution,CUDA 的虚拟 ISA,类似 LLVM IR。 | Ch8 |
| 架构机器码 | SASS | 实际硬件执行的机器码,与 SM 架构强绑定。 | Ch8 |
| 性能分析器 | Nsight Compute | NVIDIA 的 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 Profiling | CUPTI 以用户或自动 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 gap | Nsight 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 |
| NVTX | NVIDIA Tools Extension | nvtxRangePush/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 |