想象你是一家大餐厅的主厨，今晚要接待一百桌客人，每桌十道菜。如果一道道菜分别现炒现上，厨房早就乱成一锅粥。聪明的做法是：提前把菜单设计好，把食材按份切好码放整齐，在灶台边摆好随手能拿的调料，再请一台自动炒菜机专门对付最耗时的翻炒动作。这就是 GPU 上跑矩阵乘法的日常——矩阵乘法是深度学习的心脏，而 Triton 就是让你能指挥这台炒菜机的“菜谱语言”。 这篇文章会带你从零搞懂三件事：一是矩阵乘法为什么要“分块”来做，二是怎么让数据在 GPU 里流动时不卡壳（共享内存、数据排布、精度选择），三是怎么用 Triton 写出真正调用到 Tensor Core 的算子。全程用开餐厅的比喻贯穿，即使你完全没接触过 GPU 编程，也能在脑海中搭起完整的画面。

岗位关联

掌握 Triton 写 MatMul，是理解 GPU 算力上限的钥匙。无论做大模型推理优化还是训练加速，这份技能都能帮你把“能用”的算子调教成“飞快”的算子，是通往高端 AI 基础设施工程师的核心跳板。

目录
  1. 可视化演示
  2. 分块策略
  3. 共享内存
  4. 数据排布
  5. 精度体系
  6. Tensor Core
  7. Triton语法
  8. 性能调优
  9. AI 应用场景
  10. Python 演示
  11. 常见误区
  12. 面试题

可视化演示

展示Triton实现高性能矩阵乘法的核心优化：分块计算模式下可视化输出矩阵的Tile分块计算流程，精度对比模式下展示FP16与INT8在Tensor Core上的性能差异。

本节配有交互动效（进入视口后自动加载，约 3KB）。

分块计算模式 输出矩阵C=A×B被切分为Tile块逐个计算 每个Tile加载A和B的对应分块到共享内存 Cij=ΣkAik×Bkj (blocked) Triton通过分块优化GPU并行计算效率

分块策略

你是一家大餐厅的主厨，今晚要接待一百桌客人。聪明的主厨绝不会把一桌十道菜全部做完再接待下一桌，而是把菜谱拆成一道道“标准盘”：每份红烧肉、每份清蒸鱼都有固定的备料量和烹饪步骤。GPU 的处境和这位主厨一模一样——如果让每个线程单独负责整个矩阵中的一个数字，它每算一步都要重新去“大仓库”取材料，慢得让人崩溃。

矩阵乘法 C = A×B 的计算核心是：输出一个元素需要 A 中一整行和 B 中一整列。与其让每个线程单打独斗，不如把输出矩阵切成许多小方块（Tile/Block），让一组线程（Triton 里的一个 program）合力算出一个方块。这一步叫分块（Tiling）。每个输出块对应 A 的一个子块和 B 的一个子块，子块被反复使用，为后面的共享内存优化埋下伏笔。

具体地，若输出块大小为 BLOCK_M × BLOCK_N，则它需要的 A 子块是 BLOCK_M × K，B 子块是 K × BLOCK_N。分块大小是性能的命门：块太小，线程空转浪费并行度；块太大，共享内存装不下、寄存器溢出，反而更慢。

一盘红烧肉就是一个输出块，备料相当于加载 A/B 子块，翻炒相当于循环累加。把这个流程交给几百个“帮厨”（线程）并行干，整桌宴席自然能按时上齐。而分块策略就是主厨提前画好的“切菜图”——切多大、切几刀，直接决定后厨出菜效率。

C[i,j]=∑k=0K−1A[i,k]⋅B[k,j]\displaystyle C[i, j] = \sum_{k=0}^{K-1} A[i, k] \cdot B[k, j]
这个公式是说：C 矩阵第 i 行第 j 列的数字，等于 A 的第 i 行和 B 的第 j 列逐位相乘后全部加起来。
  1. 输出下标i 和 j 分别表示输出矩阵的行与列，即我们要算 C 的第 i 行第 j 列这个位置。
    C[i,j]C[i, j]
  2. 求和符号k 从 0 到 K-1 遍历，这是把 A 行和 B 列的所有对应位置都过一遍。
    ∑k=0K−1\sum_{k=0}^{K-1}
  3. 乘累加每个 k 对应一个 A 元素乘一个 B 元素，乘积全部累加进 C[i, j]。
    A[i,k]⋅B[k,j]A[i, k] \cdot B[k, j]

共享内存

主厨炒菜时，灶台边上一定要有一个“备菜台”，切好的葱姜蒜、配好的料汁都放在伸手就能拿到的地方。如果每次要用酱油都跑去仓库拿，菜早糊了。GPU 的共享内存（Shared Memory）正是这样的备菜台：它长在芯片内部，离计算单元非常近，访问速度比全局内存快一个数量级。

在分块策略之后，我们有了小块的 A 和 B，但它们还躺在全局内存这个“大仓库”里。全局内存访问延迟高达数百周期，而共享内存只有几十周期。所以正确的步骤是：先把块从全局内存搬运到共享内存，再从共享内存快速取用多次。这就是为什么 Triton 内核中 tl.load 进来的数据，编译器会自动帮助放入共享内存。

每个 BLOCK_M×BLOCK_N 的输出块，要用 A 的 BLOCK_M 行和 B 的 BLOCK_N 列各 K 次。如果不经过共享内存，每个线程都去全局内存重复取相同的数据；借助共享内存，同一份数据可以被块内所有线程重复利用——这就是数据复用（Data Reuse），也是共享内存能大幅提速的根本原因。此外，当多个线程访问共享内存的同一组 bank（存储通道）时会产生冲突，这会拖慢速度，所以高手会刻意调整数据排布来避开 bank 冲突。

灶台上的备菜台不可能放下一百桌的全部食材，只能放“这一锅”要用的。共享内存也一样，容量通常只有几十 KB，每次只能放一两个块，用完就换下一批——这就是为什么分块策略和共享内存永远搭配使用：分块决定了备菜台上摆什么，共享内存决定了摆得多顺手。

Tglobal≈200 cycles,Tshared≈30 cycles\displaystyle T_{\text{global}} \approx 200\ \text{cycles}, \quad T_{\text{shared}} \approx 30\ \text{cycles}
全局内存取数约 200 个时钟周期，共享内存约 30 个时钟周期，差了近七倍，所以“先把菜放备菜台”能带来巨大收益。
  1. 全局内存延迟访问一次全局内存大约要 200 个时钟周期，相当于从仓库跑一个来回。
    Tglobal≈200 cyclesT_{\text{global}} \approx 200\ \text{cycles}
  2. 共享内存延迟访问一次共享内存大约只要 30 个时钟周期，相当于从手边备菜台拿一下。
    Tshared≈30 cyclesT_{\text{shared}} \approx 30\ \text{cycles}
  3. 速度差距两个延迟相除约等于 6.7，这就是“把手边的菜摆好”带来的效率差距。
    TglobalTshared≈6.7\frac{T_{\text{global}}}{T_{\text{shared}}} \approx 6.7

数据排布

同样的食材，有人喜欢竖着码在冰箱里，有人习惯横着摞。取用的时候，码放方式决定了你找一根胡萝卜要翻多久。GPU 里数据在内存中的排列方式叫数据排布（Layout），它直接决定访存速度快慢——这不是玄学，而是内存硬件的物理限制。

矩阵按行主序（Row-major）存储意味着同一行的元素在内存地址上相邻；按列主序（Column-major）则同一列相邻。写矩阵乘时，如果我们用“行”的方式加载 A、用“列”的方式加载 B，内存访问会比较顺。Triton 中通过 stride（步长）参数来描述这种排布，stride 就是相邻两行（或两列）首地址之间隔了多少个元素。

GPU 读取内存的最小单位是一段连续地址（通常 32 字节），一次可以把 16 个 FP16 连续数搬进寄存器或共享内存，这叫做向量化加载。若数据不连续，比如要跨行取列，同一个加载操作就要拆成多次，效率瞬间下降。因此高手写内核时，会刻意把需要的子块组织成连续内存，或者通过转置把数据“摆顺”，让每次访存都满载而归。

把菜按“一锅菜需要的量”提前摆好，而不是散乱堆放，厨师取用时手不抖、速度快。Triton 里让 A 按行主序读取、B 按列主序读取（或做转置），本质上就是在把食材摆放成最顺手的样子——数据排布顺了，访存效率自然就上去了。

addr(i,j)=i×ld+j\displaystyle \text{addr}(i, j) = i \times \text{ld} + j
行主序下，第 i 行第 j 列元素的地址 = i 乘以行间距 ld 再加列偏移 j。同行的元素天然连续，所以扫描一行就像从冰箱第一格顺次拿到最后一格。
  1. 元素地址addr(i, j) 表示第 i 行第 j 列元素在内存中的位置，用它来定位数据。
    addr(i,j)\text{addr}(i, j)
  2. 行间距i 乘以 ld 表示跳过前面 i 整行，ld 是相邻两行首地址的距离，通常等于矩阵的列数。
    i×ldi \times \text{ld}
  3. 列偏移末尾的 j 表示在当前行内部的列偏移，因为行主序下同一行是连续存放的。
    +j+ j

精度体系

高级餐厅的食材分级是门学问：招待贵宾用 A5 和牛，做员工餐用普通牛肉，成本差上百倍，但味道也还能接受。GPU 算矩阵也有“食材等级”：FP32 是顶级和牛，FP16 是平价牛肉，INT8 是冷冻肉——各有各的适用场景，关键是别花冤枉钱。

FP32 用 32 位二进制表示一个数，精度高但占内存大、计算慢；FP16 用 16 位，显存占用减半、吞吐翻倍，但能表示的数值范围和精度都缩水；INT8 更进一步用 8 位整数，最快最省但需要量化。Tensor Core 对 FP16 和 INT8 都有专门的加速路径。大模型训练时常用混合精度（AMP），就是让计算走 FP16、主权重保留 FP32，既快又不掉精度。

FP16 的数值范围大约只有 ±65504，而 FP32 高达 ±3.4×10³⁸。若中间结果溢出，FP16 会直接变成无穷大，模型训练立刻崩溃。所以实践中要么做梯度缩放，要么保留 FP32 累加器。INT8 则需要先做量化：把浮点数值按比例缩放映射到 [-128, 127] 的整数区间，适合对精度宽容度较高的推理阶段。

全用 A5 和牛做饭不仅贵，高压锅也发挥不出顶级食材的全部优势。混合精度就像“贵菜用和牛、日常菜用牛肉”——口味过得去（精度不崩），又能把锅的使用率（算力吞吐）拉满。精度体系这门学问，就是在“味道”和“成本”之间找到最佳平衡点。

FP16:(−1)s×2e−15×(1+m/1024)\displaystyle \text{FP16}: (-1)^s \times 2^{e-15} \times (1 + m/1024)
FP16 的组成：1 位符号，5 位指数，10 位尾数。它比 FP32 省一半空间，但只有大约 3 位十进制有效数字，数值范围也小很多。
  1. 符号位 s1 个二进制位，决定这个数是正还是负。
    (−1)s(-1)^s
  2. 指数位 e5 个二进制位决定数量级，以 2 的（e-15）次方放大或缩小，所以范围大约到 ±65504。
    2e−152^{e-15}
  3. 尾数位 m10 个二进制位决定精度细节，分母 1024 来自 2 的 10 次方，这让它大约只有 3 位有效数字。
    (1+m/1024)(1 + m/1024)

Tensor Core

你终于决定买一台自动炒菜机。它不是万能的，但在“颠勺翻炒”这一件事上，效率是人工的十倍。Tensor Core 就是 NVIDIA GPU 里为矩阵乘法专门设计的“自动炒菜机”——它不做通用计算，只专精小矩阵的乘累加。

Tensor Core 是 GPU 芯片上的专用硬件单元，一次指令能在极短周期内完成一个小矩阵乘法，吞吐量比普通 CUDA 核心高出数倍。从 Ampere 架构开始，Tensor Core 支持 m16n8k16 的运算形状：M=16 行、N=8 列、K=16 深度的矩阵乘。这个操作由一个 warp（32 个线程）协同完成，每个线程持有结果矩阵的一部分片断（Fragments），最后在累加器中汇总。

程序员不需要手写底层 PTX 指令（mma.sync），Triton 的 tl.dot 就能自动把块级矩阵乘映射到 Tensor Core 上。累加器（Accumulator）的概念也很关键：矩阵乘是“乘累加”，Tensor Core 会把新乘出来的结果加到累加器的旧值上，这样多次 K 循环就能无缝衔接。

有了自动炒菜机，主厨只需把备好的菜（数据块）递进进料口、按下按钮（tl.dot），机器自己颠勺出锅。你不需要知道齿轮怎么转，只需要知道它能同时处理 16 份食材。Tensor Core 就是那个按钮背后的魔法——知道它的存在，你才能理解 Triton 为什么能跑得这么快。

D16×8=A16×16×B16×8+C16×8\displaystyle D_{16\times8} = A_{16\times16} \times B_{16\times8} + C_{16\times8}
一次 Tensor Core 运算：16×16 的 A 乘以 16×8 的 B，再加上 16×8 的累加器 C，得到 16×8 的结果 D。一步完成一个小矩阵乘加。
  1. A 的形状16 行 16 列，代表 K 方向深度为 16 的输入矩阵。
    A16×16A_{16\times16}
  2. B 的形状16 行 8 列，代表输出在 N 方向宽度为 8。
    B16×8B_{16\times8}
  3. 累加器 C16×8 的矩阵，与结果同形状，承载之前循环累积下来的部分和。
    C16×8C_{16\times8}

Triton语法

有了自动炒菜机还不算完，你得会读它的说明书。Triton 就是 GPU 的“菜谱语言”——不用你手动控制每口锅的火候和翻勺次数，只需描述性地写下“把这份食材炒熟”。它让你把注意力集中在算法本身，而不是底层硬件细节上。

传统 CUDA 编程要求你手动管理线程块、线程同步、共享内存分配和寄存器分配，繁琐且容易出错。Triton 采用“块语言”（Block Language）：你在 Python 函数中描述一个块怎么算，编译器负责把它映射到底层线程和硬件指令。这种抽象让同一个内核能自动适配不同架构的 GPU，不必为每代硬件重写一遍。

一个 Triton MatMul 内核的核心只有几行：tl.load 从全局内存加载块、tl.dot 执行块级矩阵乘并累加、tl.store 写回结果。调度循环由 for k in range(0, K, BLOCK_K) 驱动，编译器在后台自动进行共享内存分配、寄存器分配和 Tensor Core 指令映射。你写的是“想算什么”，而不是“怎么算”。

Triton 菜谱短小精悍，三行话就把“备菜、炒菜、装盘”说清楚了。编译器就是那个读菜谱的老帮厨，知道怎么把“炒熟”翻译成抽油烟机和火候的具体操作。写 Triton 的人像主厨——只关心菜怎么做，不关心煤气管道怎么走。

Cblock+=Ablock(k)×Bblock(k)\displaystyle C_{\text{block}} \mathrel{+}= A_{\text{block}}(k) \times B_{\text{block}}(k)
这是 Triton 内核里的核心计算：每个输出块在循环中不断累加 A 和 B 的子块乘积。每轮循环处理 K 方向上的 BLOCK_K 个元素。
  1. 块累加C_block 是输出块，+= 表示把每次乘出来的结果不断累加进去，这就是累积器的角色。
    Cblock+=C_{\text{block}} \mathrel{+}=
  2. A 子块切片A_block(k) 表示从全局 A 矩阵中取出当前循环步对应的 K 方向切片。
    Ablock(k)A_{\text{block}}(k)
  3. B 子块切片B_block(k) 表示对应的 B 矩阵切片，两者相乘后累加到输出块。
    Bblock(k)B_{\text{block}}(k)

性能调优

菜谱写好了，但真正的考验是高峰期能不能扛住。后厨动线——备菜台放哪、炒菜机排哪、出菜口在哪——决定了高峰时期是井然有序还是鸡飞狗跳。性能调优就是优化这条动线：先用理论算出理想时间，再用工具测量实际时间，两者差距就是优化空间。

对矩阵乘法而言，理想时间由两条路径决定：一是计算路径，总运算量除以 GPU 峰值算力；二是搬运路径，总数据量除以显存带宽。实际运行时间约等于两者中更大的那个——这就是 Roofline 模型（屋顶模型）。如果搬运时间更长，说明算子处于“带宽受限”状态，GPU 计算单元在空等数据；如果计算时间更长，则处于“计算受限”状态，已经吃满了算力。

调参技巧上，第一步是调大 BLOCK 尺寸让计算密度超过阈值，从带宽受限转向计算受限；第二步是用 NVIDIA 的 profile 工具（如 ncu）分析 L2 缓存命中率、共享内存 bank 冲突、warp 占用率，逐项排除瓶颈；第三步是关注带宽利用，尽量减少无谓的数据搬运，比如把多个小算子融合成一个内核。实际中往往要像调酱料一样反复试：BLOCK_M、BLOCK_N、BLOCK_K 三个旋钮，配合 num_stages 流水线深度，一点点逼近硬件的理论巅峰。

高峰开餐前，主厨会算一笔账：灶台火力全开每分钟出多少道菜？传菜通道一分钟能送多少道？哪个先满就是瓶颈。性能调优就是拿秒表卡出瓶颈然后对症下药——要么加火力（更大块、上 Tensor Core），要么拓宽通道（优化排布、减少搬运）。没有一招鲜，只有不停测、不停调。

t=max⁡(2MNKPpeak, MK+NK+MNBW)\displaystyle t = \max\left(\frac{2MNK}{P_{\text{peak}}},\ \frac{MK + NK + MN}{B_W}\right)
总耗时约等于“算力时间”和“搬运时间”中较大者。2MNK 是矩阵乘法的总乘加次数，MK+NK+MN 是搬运 A、B、C 三个矩阵的总数据量。
  1. 计算时间2MNK 次浮点运算除以峰值算力 P_peak，得到纯算力视角的最小时间。
    2MNKPpeak\frac{2MNK}{P_{\text{peak}}}
  2. 搬运时间MK+NK+MN 个数据元素除以显存带宽 B_W，得到纯搬运视角的最小时间。
    MK+NK+MNBW\frac{MK + NK + MN}{B_W}
  3. 取两者最大值实际瓶颈在更慢的一方，优化就应该瞄准它，而不是盲目调另一侧。
    max⁡(⋯ )\max(\cdots)

AI 应用场景

大模型推理引擎的 MatMul 优化

阶段：模型推理优化

解决的问题：LLM 推理时，注意力层和 FFN 层大量 MatMul 算子跑不满 GPU，传统 PyTorch 算子频繁启动内核、搬运开销巨大。用 Triton 重写 MatMul 并调用 Tensor Core 后，单算子吞吐可提升数倍，推理延迟显著下降。

具体案例：在 vLLM 或推理服务中，decode 阶段的 QKV 投影是典型的小批量大矩阵乘法，瓶颈往往在带宽而非算力。工程师用 Triton 写一个融合了量化反量化和 MatMul 的算子，把 KV cache 访存与计算重叠，实测首 token 延迟降低约 30%。

训练框架中的融合算子开发

阶段：训练加速

解决的问题：PyTorch 中线性层 + 激活函数 + Dropout 是三个独立内核，每个都要启动一次且来回读写显存。用 Triton 将它们融合成一个内核，减少中间张量的内存搬运。

具体案例：在 BERT 训练中，把 Linear+GeLU+Dropout 的 Roofline 瓶颈从带宽受限转为计算受限，端到端训练速度提升约 20%。训练框架工程师只需写几十行 Triton 代码，就能替换掉原本数百行 CUDA。

量化模型的 INT8 推理落地

阶段：模型压缩与部署

解决的问题：INT8 量化后，普通 CUDA 核无法利用 Tensor Core 的整数矩阵乘加速能力，需要专门算子。Triton 让 INT8 MatMul 的开发变得像写 FP16 一样简单。

具体案例：某团队将 GPT 量化模型部署到 T4 GPU，用 Triton 实现 INT8 MatMul，将量化表和缩放系数放到共享内存，QKV 投影的推理吞吐提升 2.1 倍，精度损失控制在 1% 以内。

Python 演示

python59 行
import torchimport tritonimport triton.language as tl  @triton.jitdef matmul_kernel(    a_ptr, b_ptr, c_ptr,    M, N, K,    stride_am, stride_ak,    stride_bk, stride_bn,    stride_cm, stride_cn,    BLOCK_M: tl.constexpr,    BLOCK_N: tl.constexpr,    BLOCK_K: tl.constexpr,):    pid_m = tl.program_id(0)    pid_n = tl.program_id(1)    offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)    offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)    offs_k = tl.arange(0, BLOCK_K)    a_ptrs = a_ptr + offs_m[:, None] * stride_am + offs_k[None, :] * stride_ak    b_ptrs = b_ptr + offs_k[:, None] * stride_bk + offs_n[None, :] * stride_bn    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)    for k in range(0, K, BLOCK_K):        a = tl.load(a_ptrs)        b = tl.load(b_ptrs)        acc = tl.dot(a, b, acc)        a_ptrs += BLOCK_K * stride_ak        b_ptrs += BLOCK_K * stride_bk    offs_cm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)    offs_cn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)    c_ptrs = c_ptr + offs_cm[:, None] * stride_cm + offs_cn[None, :] * stride_cn    tl.store(c_ptrs, acc)  def matmul(a, b):    M, K = a.shape    K, N = b.shape    c = torch.empty((M, N), device=a.device, dtype=torch.float32)    BLOCK_M, BLOCK_N, BLOCK_K = 64, 64, 32    grid = (M // BLOCK_M, N // BLOCK_N)    matmul_kernel[grid](        a, b, c, M, N, K,        a.stride(0), a.stride(1),        b.stride(0), b.stride(1),        c.stride(0), c.stride(1),        BLOCK_M=BLOCK_M, BLOCK_N=BLOCK_N, BLOCK_K=BLOCK_K,    )    return c  if __name__ == "__main__":    torch.manual_seed(0)    a = torch.randn(256, 256, device="cuda")    b = torch.randn(256, 256, device="cuda")    c = matmul(a, b)    ref = a @ b    print("max error:", (c - ref).abs().max().item())
目的：用 Triton 从零实现一个分块矩阵乘法内核，并通过 tl.dot 自动调用 Tensor Core 加速，最后与 PyTorch 结果对比验证正确性。
  1. pid_m = tl.program_id(0)拿到当前程序在 M 方向的块编号，相当于知道自己是负责哪一排宴席的帮厨。
  2. offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)算出输出块的行索引向量，对应本块要服务的 64 桌客人。
  3. tl.load(a_ptrs)从全局内存加载 A 子块到高速暂存区，相当于把这道菜需要的食材搬到备菜台。
  4. acc = tl.dot(a, b, acc)一次块级矩阵乘并累加，Triton 编译器自动映射到 Tensor Core，是全场最核心的一勺翻炒。
  5. tl.store(c_ptrs, acc)把算好的输出块写回全局内存，相当于把做好的菜端到出菜口。

常见误区

误区“BLOCK 尺寸越大，每次算的越多，肯定越快。我直接把 BLOCK 设成 256×256，一步到位。”

正解分块大有两座大山压着：一是共享内存容量有限，块太大直接装不下或换入换出频繁；二是块太大导致块数量太少，GPU 上并行执行的线程组不够，大量计算单元空闲。块大小必须配合硬件参数，通常 BLOCK_M/N 取 64~128、BLOCK_K 取 32~64 是安全起点，再根据 profile 结果微调。

为何常见：“一次多做点”符合日常直觉，但忽略了 GPU 是“人多力量大”的并行机器，块多才是充分利用算力的关键。

误区“FP16 只有 3 位有效数字，训练和推理用了肯定大幅掉点，还是老老实实用 FP32。”

正解FP16 的精度损失在大多数场景下可以忽略，因为 Triton 的 tl.dot 支持 FP32 累加器——中间累积用高精度，只有输入是半精度。再配合混合精度技术（主权重 FP32、计算 FP16）和损失缩放，模型精度几乎无损，而吞吐直接翻倍。INT8 才需要更谨慎的量化校准。

为何常见：人们把“输入精度”和“计算精度”混为一谈，忽略了累加器可以更宽这个关键设计。

误区“Triton 写起来是方便，但它是高级语言，编译器生成的代码肯定不如我手写的 CUDA 快。”

正解Triton 编译器的块级抽象让它可以全局优化内存布局、寄存器分配和指令调度，在 MatMul 这类规整算子上，性能接近甚至超越手写 CUDA。它能自动适配不同 GPU 架构的 Tensor Core 指令，而手写 CUDA 换个架构可能还要重写。真正复杂的自定义算子才需要 CUDA 的精细控制。

为何常见：很多人有“手工的一定比自动的好”的直觉，忽略了编译器几十年的优化积累和硬件架构的快速迭代。

误区“我把 BLOCK_M/N/K 调成文档推荐的 64/64/32，性能应该就到顶了。”

正解BLOCK 只是第一层。数据排布（行主序还是列主序）、共享内存 bank 冲突、num_stages 流水线深度、内核启动开销、与前后算子的融合，每个环节都可能成为新瓶颈。Roofline 模型告诉你先判断是计算受限还是带宽受限，再针对性下手。调参是系统工程，不是拧一个旋钮。

为何常见：把多维优化问题简化成一维，是新手最容易犯的归因错误，调完发现还有两倍差距却不知从何下手。

面试题

Q1. 在矩阵乘法 C = A × B 中，C[i][j] 的计算方式是？

单选题 · easy

  • A.A 第 i 行与 B 第 j 列的点积← 正确答案
  • B.A 第 j 行与 B 第 i 列的点积
  • C.A 第 i 列与 B 第 j 行的点积
  • D.A 第 i 行与 B 第 j 行的点积

答案：A。C[i][j] 等于 A 的第 i 行与 B 的第 j 列逐元素相乘再累加，这正是点积的定义。其余选项把行列关系搞反了，就像做菜把食材顺序颠倒了，味道自然不对。

Q2. 矩阵乘法为什么要分块（Tiling）？

单选题 · mid

  • A.减少 GPU 线程数量
  • B.复用数据减少慢速访存← 正确答案
  • C.降低计算精度损失
  • D.简化编译器开发难度

答案：B。分块让一个块内的数据在共享内存中被反复复用，避免每个线程都去全局内存重复取数。全局内存慢、共享内存快，分块就是把“频繁去仓库取菜”改成“一次搬一筐到备菜台”。

Q3. 关于共享内存，以下哪些说法正确？

多选题 · mid

  • A.比全局内存访问快得多← 正确答案
  • B.容量远大于全局内存
  • C.是芯片内高速暂存区← 正确答案
  • D.数据断电后仍不丢失

答案：AC。共享内存在 GPU 芯片内部，延迟约 30 周期，远快于全局内存的 200 周期，但容量只有几十 KB。它是临时暂存区，断电数据即失。备菜台虽快，但放不下整个仓库的货。

Q4. FP16 相比 FP32 的主要缺点是？

单选题 · mid

  • A.计算速度反而更慢
  • B.数值范围和精度变小← 正确答案
  • C.占用显存空间更多
  • D.无法执行乘法运算

答案：B。FP16 只有 5 位指数和 10 位尾数，范围约 ±65504，精度约 3 位有效数字，都比 FP32 小。但因为位宽减半，它更快更省显存。就像平价牛肉便宜易熟，但口感和等级不及和牛。

Q5. Tensor Core 的主要用途是什么？

单选题 · easy

  • A.通用逻辑运算
  • B.矩阵乘加计算← 正确答案
  • C.全局内存分配
  • D.图形渲染处理

答案：B。Tensor Core 专精小矩阵的乘累加（如 16×8×16），吞吐量远超通用 CUDA 核，是 GPU 上专门用来“颠勺翻炒”的自动炒菜机。

Q6. 关于 Triton 的 tl.dot，以下哪些说法正确？

多选题 · hard

  • A.接受块级张量作为参数← 正确答案
  • B.编译器自动映射到 Tensor Core← 正确答案
  • C.必须手写 PTX 底层指令
  • D.只能用于 INT8 数据类型

答案：AB。tl.dot(a, b, acc) 接收三个块级张量，编译器自动生成底层 mma 指令映射到 Tensor Core。程序员无需手写 PTX，精度由输入数据类型决定，支持 FP16/INT8 等多种类型。

Q7. 行主序存储时，A[i][j] 的内存地址计算公式是？

单选题 · mid

  • A.i*ld + j 地址公式← 正确答案
  • B.j*ld + i 地址公式
  • C.i + j 地址公式
  • D.i*j 地址公式

答案：A。行主序下同一行连续存储，地址 = 行号 i 乘以行间距 ld，再加上列偏移 j。就像按层摆放食材，先找到第 i 层，再在该层找第 j 格。

Q8. Roofline 模型的核心结论是什么？

单选题 · hard

  • A.瓶颈取决于算力与带宽← 正确答案
  • B.块越大性能就一定越好
  • C.精度越高性能就一定越好
  • D.共享内存越大性能越好

答案：A。实际运行时间约等于“算力视角时间”和“带宽视角时间”中较大者。优化应瞄准瓶颈那一侧：算力受限就加大块用 Tensor Core，带宽受限就优化排布减少搬运。

Q9. 使用 Triton 编程的主要好处有哪些？

多选题 · easy

  • A.无需手动分配共享内存← 正确答案
  • B.用 Python 风格编写内核← 正确答案
  • C.编译器自动处理依赖调度← 正确答案
  • D.完全无需理解 GPU 硬件

答案：ABC。Triton 自动管理共享内存和寄存器，Python 风格描述块计算，编译器负责依赖分析和指令映射。但写高性能算子仍需理解 GPU 的访存和并行模型，D 选项太过绝对。

Q10. Ampere 架构 Tensor Core 的 m16n8k16 指令形状含义是？

单选题 · extreme

  • A.一次算 16 行 8 列深 16← 正确答案
  • B.一次算 8 行 16 列深 16
  • C.一次算 16 行 16 列深 8
  • D.一次算出整个大矩阵

答案：A。m16n8k16 指 M=16、N=8、K=16，由 warp 中 32 个线程协作完成一次小矩阵乘加。它像自动炒菜机一次最多炒 16 份菜、每份 8 个配菜、深度 16 层，是硬件层面的指令形状。