想象你是一家大型食堂的后厨主厨。过去你只有一口小炒锅，一次只能炒一粒米，几千人排队饿得眼冒金星。后来后厨搬来一口巨型铁锅——只要把食材整整齐齐码成方阵倒进去，一锅就能炒出几百粒米，火候均匀、香气四溢。这口“大铁锅”，就是 GPU 里的 Tensor Core。它专为矩阵乘法而造，一次指令同时完成一个子矩阵的乘法和加法，把 AI 训练和推理中最常见的“矩阵乘加”动作，从一颗颗单炒升级成整锅快炒。但大铁锅也有讲究：食材切得太碎，锅空着一大半；食材切得太大，锅口塞不下。利用率高低，决定了这口锅是满负荷爆炒，还是大半时间在空烧等食材。本文就用“后厨炒菜”的视角，带你从概念、指标、策略到实战，把 Tensor Core 利用率优化这件事彻底讲透。

岗位关联

对 AI 从业者而言，Tensor Core 利用率优化直接决定了大模型训练和推理的速度与成本。掌握它，你就能用更少的 GPU 完成更大的模型，这是从“能跑”到“跑得快、跑得省”的关键一跃。

目录
  1. 可视化演示
  2. 基础概念
  3. 影响因子
  4. 性能指标
  5. 优化策略
  6. 存储访问
  7. 案例调优
  8. AI 应用场景
  9. Python 演示
  10. 常见误区
  11. 面试题

可视化演示

动态展示Tensor Core矩阵乘加运算、维度对齐策略、算子融合与异步拷贝流水线四个视角，直观理解利用率优化核心机制。

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

基础概念

假设食堂后厨以前只有小炒锅，一颗米一颗米地炒，一个饭点最多服务几十人。后来搬来一口巨大的铁锅——只要把米和菜整齐码成“方阵”倒进去，一次就能同时炒出几百粒米。Tensor Core 就是 GPU 里这口专为矩阵乘法定制的大铁锅：它不再像普通 CUDA 核心那样一次算一个数，而是把一个“子矩阵”整个丢进去，一次完成一个矩阵块的全部乘法和加法。这种设计，让 AI 中最常见的“大矩阵相乘”从效率极低的逐个相乘，变成了硬件级的整块爆炒。

从技术上讲，Tensor Core 是 NVIDIA 从 Volta 架构开始在每块 GPU 中内置的专用矩阵乘加单元。它执行的是一条融合指令：把矩阵 A 与矩阵 B 相乘，再把结果加到矩阵 C 上，最终输出 D。这里的“融合”就像炒菜时“翻炒的同时顺手加调料”——乘法和加法在一次硬件操作中完成，省去来回搬运中间结果的麻烦。Tensor Core 支持 FP16、BF16、INT8 等低精度数据类型，配合混合精度技术，峰值吞吐往往是普通 FP32 CUDA 核心的数倍甚至一个数量级。

回到食堂的比喻：大铁锅一次能炒的方阵越大，单位时间内出锅的饭菜就越多。但铁锅搬来只是第一步，食材怎么切、厨师怎么配合，才是影响这口锅能不能真正“满负荷爆炒”的关键。

Dm×n=Am×k×Bk×n+Cm×n\displaystyle D_{m\times n} = A_{m\times k} \times B_{k\times n} + C_{m\times n}
这个公式说明：Tensor Core 的核心本领是“先乘后加”，把输出矩阵 D 一次性算出来。示意图里的 m、n、k 分别是行数、列数和内部分量，它们共同决定了一次铁锅翻炒的“食材总量”。
  1. 矩阵维度检查乘法的前提是 A 的内维度 k 与 B 的内维度 k 相等，这样 A 的一行才能与 B 的一列逐点相乘再求和。
    Am×k, Bk×nA_{m\times k},\,B_{k\times n}
  2. 乘加融合操作D 的每一个元素等于 A 对应行与 B 对应列的点积，然后再加上 C 的对应元素，这就是“乘加”的含义。
    Dij=∑r=1kAirBrj+CijD_{ij} = \sum_{r=1}^{k} A_{ir}B_{rj} + C_{ij}
  3. Tensor Core 硬件执行Tensor Core 不是按单个元素算，而是把 16×16×16 的一整块数据同时送入运算单元，一条指令完成大量 D_{ij}。
    Tensor Core 单条指令处理 16×16×16\text{Tensor Core 单条指令处理 }16\times16\times16

影响因子

铁锅炒菜，火候旺不旺，不光看锅有多大，还要看食材怎么切、几个厨师怎么配合。Tensor Core 的世界里，食材就是矩阵的“分块”（Tile），厨师就是 warp——一组 32 个线程。而每一次“颠勺”动作，就是一条 HMMA 指令。如果食材切得和锅口不匹配，或者厨师们等待口令停太久，锅就算再大也会空烧。

HMMA 全称是 Half-precision Matrix Multiply Accumulate，是 Tensor Core 执行矩阵乘加的核心指令。以最常见的 HMMA.16816 为例，它让一个 warp（32 个线程）协同完成一个 16×16×16 的矩阵块运算。也就是说，Tensor Core 的调度粒度是一个 warp，而不是单个线程。这种“一个 warp 一条指令炒一锅”的模式，决定了利用率的高低往往取决于：tile 尺寸是否对齐到 16 的倍数、warp 调度之间有没有空隙、以及寄存器是否足够放下当前锅里的菜。

如果 tile 是 17×17，就会有一行一列“超出锅口”，白白浪费部分算力；如果寄存器不够用，数据会被挤到本地内存，相当于把菜端出备菜台再端回来，铁锅只能干等。因此，在写矩阵乘 kernel 时，必须精挑细选 tile 形状，让每次 HMMA 都尽量把整口锅装满。

所以，影响因子可以概括为三件事：食材大小（tile）、厨师动作（HMMA 指令）和调度节奏（warp 是否连续干活）。这些因素最终都会被一个统一的指标收拢——那就是“利用率”。

nHMMA=⌈M16⌉×⌈N16⌉×⌈K16⌉\displaystyle n_{\text{HMMA}} = \left\lceil \frac{M}{16} \right\rceil \times \left\lceil \frac{N}{16} \right\rceil \times \left\lceil \frac{K}{16} \right\rceil
这个公式告诉我们：处理一个 M×N×K 的矩阵乘，需要多少条 HMMA 指令。所谓"16"就是铁锅锅口的长宽深——一次最多装 16×16×16 个元素。
  1. 按锅口尺寸划分HMMA 固定一次处理 16×16 的输出块，所以先把矩阵在 M 和 N 维度各切成 16 大小的块。
    ⌈M16⌉,⌈N16⌉\left\lceil \frac{M}{16} \right\rceil,\quad \left\lceil \frac{N}{16} \right\rceil
  2. K 维也要分块K 是求和方向，也需要切成 16 长度的段，每段都要参与一次矩阵乘累加。
    ⌈K16⌉\left\lceil \frac{K}{16} \right\rceil
  3. 总指令数相乘输出块数乘以 K 维分段数，才是完成整个矩阵乘需要的 HMMA 指令总数。
    nHMMA=输出块数×K 维段数n_{\text{HMMA}} = \text{输出块数}\times \text{K 维段数}

性能指标

后厨要看一口锅好不好用，不能只看锅有多大，得看它“真正冒火翻炒的时间”占了多少。Tensor Core 利用率就是这样一个“火候指标”：它等于铁锅真正计算的时间占总时间的比例。如果一段时间里锅有 70% 的时间在炒菜、30% 的时间在等食材，那利用率就是 70%。

在 GPU 性能分析中，利用率是衡量 Tensor Core 是否被充分使用的最直观数字。它由两部分决定：一部分是有效计算时间 T_compute，也就是 Tensor Core 正在执行 HMMA 指令的时间；另一部分是空闲等待时间 T_idle，包括等数据从全局内存搬回、等 warp 之间的同步、等寄存器依赖解决等。占空比（duty cycle）也是类似概念，描述一个调度周期里实际干活的比例。

训练和推理场景往往差别很大：训练时 batch size 很大，矩阵维度饱满，Tensor Core 容易打满；推理时 batch 常常只有 1，矩阵变成“细长条”，锅口装不满，利用率就会急剧下降。因此，在评估模型性能时，不能只看理论峰值 TFLOPS，更要看“实际效率占峰值的比例”，即理论峰值比。

一旦知道利用率低，问题就变成了：怎么让 T_idle 趋近于零？答案是让食材一波接一波下锅，不留给铁锅空烧的机会——这正是下一节要讲的分块和双缓冲。

U=TcomputeTcompute+Tidle\displaystyle U = \frac{T_{\text{compute}}}{T_{\text{compute}} + T_{\text{idle}}}
利用率 = 干活时间 /（干活时间 + 发呆时间）。这个值越接近 1，说明铁锅越是在满负荷爆炒；越接近 0，说明锅大部分时间在空烧。
  1. 计算时间Tensor Core 真正执行矩阵乘加指令的时间，也就是铁锅烧红翻炒的时间。
    TcomputeT_{\text{compute}}
  2. 等待时间数据没到位、同步没完成、寄存器依赖未解决造成的空闲时间，相当于厨师等菜下锅。
    TidleT_{\text{idle}}
  3. 利用率汇总两者相加是总时间，计算时间占比就是利用率，取值范围在 0 到 1 之间。
    U=TcomputeTcompute+TidleU = \frac{T_{\text{compute}}}{T_{\text{compute}} + T_{\text{idle}}}

优化策略

把大订单切成能塞进锅口的食材方阵，同时一边炒菜一边把下一份食材提前端到锅边——这就是 Tensor Core 优化里最经典的两招：分块（Tiling）和双缓冲（Double Buffering）。分块让数据在共享内存和寄存器中被反复利用，减少从“冰箱”（全局内存）取菜的次数；双缓冲则让取菜和炒菜同时进行，铁锅不再因为等菜而熄火。

在 CUDA 编程中，使用 WMMA（Warp Matrix Multiply Accumulate）API 可以直接编写 Tensor Core 矩阵乘。它把程序员从繁琐的 PTX 汇编中解放出来，同时允许指定 tile 大小。比如 16×16×16 或 32×32×16 的分块，正好对应 HMMA 指令的支持范围。除了调用 API，优化者还需要手动做分块计算：把输出矩阵 M×N 划分成若干个 T×T 的小块，每个块由一组线程负责，循环累加 K 维的数据。

双缓冲机制是进一步提高利用率的关键：在共享内存中准备两个缓冲区，一个用于当前 tile 的计算，另一个用于预取下一个 tile 的数据。这样，全局内存的数据搬运可以和 Tensor Core 的矩阵运算完全重叠。同时，还要尽量减少 __syncthreads() 同步点——同步就像厨师们必须停下锅等所有人喊“一二三”才开炒，同步越多，锅空烧越久。

分块和双缓冲让食材源源不断下锅，但还有一个隐藏瓶颈在等着：即便火候再旺，如果传送带送菜太慢、备菜台格子互相打架，铁锅照样喂不饱。下一节我们看看后厨的物流系统——存储访问。

B=⌈MT⌉×⌈NT⌉\displaystyle B = \left\lceil \frac{M}{T} \right\rceil \times \left\lceil \frac{N}{T} \right\rceil
分块数 = 行方向块数 × 列方向块数。T 是每块边长，T 越大块越少，块间同步开销越小；T 太大又可能超过铁锅和盘子的容量，造成溢出。
  1. 选择 tile 边长 TT 通常取 16、32、64 等 16 的倍数，保证每次 HMMA 都装满锅口。
    T∈{16,32,64,… }T \in \{16, 32, 64, \dots\}
  2. 计算行列方向块数矩阵在行方向被切成 ⌈M/T⌉ 块，在列方向被切成 ⌈N/T⌉ 块。
    ⌈MT⌉,⌈NT⌉\left\lceil \frac{M}{T} \right\rceil,\quad\left\lceil \frac{N}{T} \right\rceil
  3. 总块数等于两者乘积每个块交给一组线程处理，块数是调度和并行的基础单位。
    B=⌈MT⌉×⌈NT⌉B = \left\lceil \frac{M}{T} \right\rceil \times \left\lceil \frac{N}{T} \right\rceil

存储访问

就算铁锅再大、火候再旺，如果食材在“冰箱—备菜台—铁锅”这条路上卡住了，锅也只能干烧。在 GPU 里，这条路线就是全局内存→共享内存→寄存器。Tensor Core 的利用率，说到底取决于数据能不能按时送到。后厨中最常见的两个问题，一是送货太零散（内存访问不合并），二是多个厨师抢同一个备菜台格子（bank 冲突）。

优化全局内存访问时，最关键的是让同一个 warp 里的 32 个线程访问连续的地址，这叫“内存合并访问”（coalesced memory access）。如果线程各访各的、地址跳来跳去，一次内存请求就会被拆成几十次，相当于传送带一趟只运一粒米，效率极低。共享内存则被分成多个 bank（格子），如果多个线程同时访问同一个 bank，就会发生冲突，硬件只能逐个服务。

内存对齐和向量化加载也是常见手段：确保 tile 数据在内存中的起始地址对齐到 16 字节，并用 float4 等向量类型一次性读取 4 个数据，相当于把四个小包裹捆成一盒搬。为了定存储优化方向，我们可以用一个核心指标——算术强度（Arithmetic Intensity）来衡量：单位字节访存能换来多少计算。算术强度越高，说明铁锅每等来一批食材都能炒很久，利用率越容易拉高。

后厨的物流理顺之后，铁锅终于能持续冒香气了。但每个菜谱的实际情况都不一样，怎么科学地找到瓶颈？这时候就需要搬出监控探头——性能分析工具。

I=2mnkBytesA+BytesB+BytesC\displaystyle I = \frac{2mnk}{\text{Bytes}_A + \text{Bytes}_B + \text{Bytes}_C}
算术强度 = 浮点运算量 / 访存字节数。它像“每斤食材能炒出几道菜”的指标；强度越高，就越不依赖送菜速度，Tensor Core 越容易跑满。
  1. 计算运算量矩阵乘 A_{m×k} × B_{k×n} 需要大约 2mnk 次浮点运算（乘加各算一次）。
    FLOPs≈2mnk\text{FLOPs} \approx 2mnk
  2. 统计访存量从全局内存读取 A、B，写回 C，共需要三类矩阵的字节数之和，精度越高字节越多。
    BytesA+BytesB+BytesC\text{Bytes}_A + \text{Bytes}_B + \text{Bytes}_C
  3. 两者相除得到强度算术强度越高，单位访存支撑的运算越多，Tensor Core 越不容易因等数据而空烧。
    I=FLOPsBytesI = \frac{\text{FLOPs}}{\text{Bytes}}

案例调优

一个真实的后厨，不能只靠感觉判断火候，得看监控、看数据。GPU 优化也一样：先用 Nsight 之类的性能分析工具“偷看后厨”，看看铁锅到底在炒菜还是在空烧；再针对空烧原因调整菜谱；最后用加速比验证每一步是否真的有效。这就是 Tensor Core 利用率优化的完整闭环。

假设我们要优化一个 Transformer 模型中的全连接层。先用 Nsight Compute 跑一遍 profiling，看到 SM 占用率很高，但 Tensor Active 比率只有 40%——这意味着锅在灶上放着，却没怎么开火。进一步分析发现，kernel 里 tile 只设成 8×8，小于 HMMA 要求的 16×16，每一条指令都有一半锅口空着。于是做三件事：把 tile 改成 16×16 或 32×32，使用 WMMA API 显式调用 Tensor Core，把数据切成 FP16 并配合混合精度。修改后再次 profiling，Tensor Active 比率提升到 80%。

调优的效果用加速比来衡量：优化前的总时间除以优化后的总时间。如果 S 大于 1，说明铁锅这顿饭做得更快。在实际项目中，这个过程需要反复迭代：跑一次分析、改一版代码、再测一次性能。特别是面对 Attention 里的 QK^T 矩阵乘时，它的形状是 [B, H, S, S]，S 往往不是 16 的倍数，需要做形状填充或将多个小矩阵拼接成大矩阵，才能让锅口装满。

从拿到一口新铁锅，到让它每一秒都爆炒出最大火力，靠的正是“概念→指标→策略→测量”这套循环。铁锅不会自己变快，但懂得看火候、调食材的厨师，能让同样一口锅发挥出翻倍的速度。

S=TbeforeTafter\displaystyle S = \frac{T_{\text{before}}}{T_{\text{after}}}
加速比 = 优化前耗时 / 优化后耗时。这个比值像“这次改菜谱后上菜快了几倍”；比值越大，说明优化越成功。
  1. 记录优化前基线时间在改代码之前，先完整跑一遍 kernel，记录耗时 T_{\text{before}}，作为对比基准。
    TbeforeT_{\text{before}}
  2. 记录优化后时间应用 tile 调整、WMMA、双缓冲等策略后，重新测量耗时 T_{\text{after}}。
    TafterT_{\text{after}}
  3. 计算加速比两者相除就是性能提升倍数，S>1 表示优化有效，S<1 则说明改坏了。
    S=TbeforeTafterS = \frac{T_{\text{before}}}{T_{\text{after}}}

AI 应用场景

大模型训练加速

阶段：模型训练与性能调优

解决的问题：训练 Llama 这类大模型时，全连接层和 Attention 层的矩阵乘占 90% 以上算力。传统直接调用通用 GEMM 库可能因 tile 形状不匹配而无法打满 Tensor Core，导致训练效率低、GPU 成本高。通过手动调整 tile 大小、使用 WMMA API 和混合精度，可以把 Tensor Core 利用率从 50% 提升到 80% 以上。

具体案例：某团队训练 7B 参数模型时，发现 FFN 层的 K 维为 4096，恰好是 16 的倍数，但默认库把 tile 切成 8×8。于是重写为 32×32×16 的 Tensor Core kernel，并用双缓冲隐藏数据搬运，训练吞吐提升约 1.7 倍，单卡训练时间从 42 天降到 25 天。

大模型推理延迟优化

阶段：模型推理部署

解决的问题：在线推理时 batch size 往往为 1，矩阵变成“细长条”，Tensor Core 锅口装不满，利用率极低。传统静态 batching 无法解决。利用动态 batch 把多个请求拼成大矩阵，同时将 KV Cache 分块并融合 Attention 中的矩阵乘，可以让单次请求的延迟下降。

具体案例：某对话机器人服务使用 FP16 推理，单 batch 时 Tensor Core 利用率只有 25%。通过动态 batch 到 8 个请求，并将 QK^T 和 PV 两个矩阵乘合并为更饱满的形状，端到端吞吐提升 3.1 倍，P95 延迟反而下降 40%。

边缘设备低精度推理

阶段：模型量化与边缘部署

解决的问题：边缘 GPU 算力有限，带宽更有限。单纯使用 FP16 可能仍然带宽受限，利用率上不去。将权重和激活量化为 INT8，利用 Tensor Core 对 INT8 的支持，配合内存对齐和向量化访存，可以在相同功耗下获得 2-4 倍加速。

具体案例：自动驾驶平台部署目标检测模型，FP16 时由于内存带宽不足，Tensor Core 利用率仅 30%。改为 INT8 量化并做通道对齐，算术强度提升约 2 倍，利用率升至 70%，推理时延从 35ms 降到 12ms。

Python 演示

python34 行
import math def estimate_utilization(M, N, K, tile, peak_tflops, bandwidth_gbps, dtype_bytes=2):    # 计算 Tensor Core 理论上需要多少 FLOPs    flops = 2 * M * N * K    # 如果整口锅全速运转，需要多少秒    compute_time = flops / (peak_tflops * 1e12)     # 估算从全局内存搬运 A、B、C 需要的总字节数    bytes_total = (M * K + K * N + M * N) * dtype_bytes    memory_time = bytes_total / (bandwidth_gbps * 1e9)     # 分块数越多，同步等待开销越大，模拟 tile 过小导致的空烧    num_blocks = math.ceil(M / tile) * math.ceil(N / tile)    sync_overhead = num_blocks * 2e-6     # 总时间取计算与访存的大者，再加上同步等待    total_time = max(compute_time, memory_time) + sync_overhead    util = compute_time / total_time     return {        "flops": flops,        "compute_time": compute_time,        "memory_time": memory_time,        "num_blocks": num_blocks,        "total_time": total_time,        "utilization": util,    } if __name__ == "__main__":    for tile in [8, 16, 32, 64]:        r = estimate_utilization(1024, 1024, 1024, tile, 312, 2000, 2)        print(f"tile={tile:3d} blocks={r['num_blocks']:4d} "              f"util={r['utilization'] * 100:.1f}%")
目的：这个“后厨火候计算器”用 Python 模拟 Tensor Core 利用率的估算过程。输入矩阵维度、tile 大小、理论峰值算力和带宽，输出预计利用率。运行后你会看到：tile 从 8 增大到 64，分块数减少，同步开销下降，利用率逐步上升——直观展示分块对利用率的影响。
  1. flops = 2 * M * N * K矩阵乘的浮点运算量，相当于这一锅菜总共要炒多少次。
  2. compute_time = flops / (peak_tflops * 1e12)一口全速铁锅把菜炒熟需要的最短时间。
  3. bytes_total = (M * K + K * N + M * N) * dtype_bytes从冰箱搬 A、B、C 三个大冰箱格子的总送货量。
  4. num_blocks = math.ceil(M / tile) * math.ceil(N / tile)食材被切成的块数，块太多会导致频繁放下锅等口令。
  5. sync_overhead = num_blocks * 2e-6估算每块菜下锅前厨师们互相等待的时间。
  6. util = compute_time / total_time最终利用率，等于真正炒菜时间占总时间的比例。

常见误区

误区既然小 tile 会让 Tensor Core 空烧，那把 tile 调到 128、256 是不是更好？越大越能装满锅，利用率一定越高。

正解tile 过大会超过共享内存和寄存器的容量，数据被迫溢出到本地内存。这就像食材太多，厨师手边的盘子根本放不下，只能把菜端回备菜台再取，反而让铁锅长时间空等。tile 需要和硬件资源匹配，常用优化区间是 16×16 到 64×64，而不是无脑加大。

为何常见：初学者只看到“大块减少同步开销”的一面，忽略了寄存器溢出和共享内存容量限制这些反面代价。

误区我在代码里没看到 wmma::mma_sync，所以 Tensor Core 没用上，利用率低；只要显式用 WMMA API，利用率一定高。

正解cuBLAS、cuDNN 等库内部可能已经在调用 Tensor Core，但默认 kernel 的参数不一定适合你的矩阵形状。即使显式调用 WMMA，如果 tile 尺寸不匹配、内存访问不合并、warp 调度有空洞，利用率照样上不去。API 只是上车方式，车能不能跑得快还取决于路况和司机。

为何常见：很多人把“没有手写 WMMA”等同于“没使用 Tensor Core”，忽略了深度学习框架底层已经做了大量封装。

误区我的模型是 FP32 权重，Tensor Core 只支持 FP16，所以我没法享受加速。

正解现代 GPU 的 Tensor Core 还支持 TF32、BF16、INT8 等格式。TF32 可以把 FP32 数据直接送入 Tensor Core，而混合精度训练则通过 master weight 保持 FP32，前向和反向计算用 FP16/BF16。就算不想改代码，cuBLAS 也会用 TF32 路径自动加速 FP32 矩阵乘。

为何常见：大家熟知的“Tensor Core 需要 FP16”来自早期 Volta 宣传，忽略了架构迭代后进出的 TF32、BF16 等新格式。

误区cuBLAS 是 NVIDIA 官方库，肯定已经把 Tensor Core 优化到极致了，我直接用就行，不用管 tile 和内存。

正解cuBLAS 要覆盖所有可能的 shape，所以它的 heuristics 不一定选中最适合你场景的 kernel。当 M、N、K 是特殊值（比如 77、192 这种非 16 倍数）时，手动定制 kernel 或调整分块、对齐方式，往往会比默认 cuBLAS 快 20%-50%。

为何常见：官方库的“通用最优”被误认为“所有场景最优”，实际 GPU 优化高度依赖具体形状和硬件代际。

误区只要 Tensor Core 利用率能跑到 95% 以上，这个 kernel 就绝对最优了，不需要再看别的。

正解Tensor Core 利用率只是手段，端到端延迟和功耗才是最终目标。有时为了降低内存带宽压力或让其他模块并行，会有意降低利用率换取整体吞吐提升。比如一个小矩阵如果用 Tensor Core 跑，启动开销占比太高，反而不如合并到相邻算子里更划算。

为何常见：性能优化者容易陷入单指标崇拜，忘了木桶效应：利用率只是那块最显眼的木板，不是整只桶。

面试题

Q1. Tensor Core 最适合加速下面哪种计算？

单选题 · easy

  • A.矩阵乘加← 正确答案
  • B.向量加法
  • C.标量取模
  • D.像素滤波

答案：A。Tensor Core 是专为矩阵乘加 D=A×B+C 设计的硬件单元，像大铁锅适合一次炒整块菜；向量加法、标量取模这类操作走普通 CUDA 核心就够，用大铁锅反而浪费。

Q2. Tensor Core 执行乘加融合运算的公式是？

单选题 · easy

  • A.D=A×B+C← 正确答案
  • B.D=A+B+C
  • C.D=A×B×C
  • D.D=A−B×C

答案：A。Tensor Core 的核心能力是“先乘后加”的融合操作 D=A×B+C。这就像炒菜时翻炒和加调料一步完成，省去了中间取放食材的额外动作。

Q3. 一条 HMMA.16816 指令通常由多少个线程协作完成？

单选题 · easy

  • A.32 个线程← 正确答案
  • B.16 个线程
  • C.64 个线程
  • D.8 个线程

答案：A。一个 warp 正好是 32 个线程，是 GPU 调度与执行的基本单位。HMMA.16816 由整个 warp 协作完成，就像四位厨师同时抬一口大锅，而不是一个人单干。

Q4. 估算 HMMA 指令数时，最关键的影响因素是什么？

单选题 · mid

  • A.矩阵维度能否被 16 整除← 正确答案
  • B.GPU 主频高低
  • C.编译器版本新旧
  • D.线程块数量多少

答案：A。HMMA 每次处理 16×16×16 的块，所以指令数 = ⌈M/16⌉×⌈N/16⌉×⌈K/16⌉。矩阵维度与 16 的整除关系决定了锅口是否装得满，比主频、编译器、线程块数量更本质。

Q5. 利用率公式 U = T_compute / (T_compute + T_idle) 中，T_idle 代表？

单选题 · mid

  • A.Tensor Core 空等的时间← 正确答案
  • B.GPU 整卡关机的时间
  • C.编译器优化的时间
  • D.数据传输的全部时间

答案：A。T_idle 是 Tensor Core 因数据未到、同步未完成等原因而空烧的时间。它不等同于所有数据传输时间，因为双缓冲可以让一部分传输与计算重叠，不产生等待。

Q6. 哪些措施能够有效提升 Tensor Core 利用率？

多选题 · mid

  • A.使用双缓冲隐藏访存延迟← 正确答案
  • B.将 Tile 尺寸调成 16 的倍数← 正确答案
  • C.关闭共享内存改用全局内存
  • D.让同一 warp 的访存地址连续合并← 正确答案

答案：ABD。双缓冲让搬运和计算重叠，Tile 对齐 16 保证 HMMA 锅口装满，合并访存减少传送带宽压力。关闭共享内存会切断备菜台直接导致数据搬运变慢，反而降低利用率。

Q7. 矩阵乘的算术强度高意味着什么？

单选题 · hard

  • A.每单位访存能执行更多运算← 正确答案
  • B.访存次数一定更多
  • C.利用率一定会更低
  • D.不需要共享内存了

答案：A。算术强度 = FLOPs / Bytes，强度高说明每次从内存取数据都能支撑大量计算，对 Tensor Core 这种计算快的锅来说更容易跑满。访存次数多反而可能是强度低的表现。

Q8. Nsight 分析中发现 Tensor Core 长期处于活跃状态，但 HMMA 指令发射速率偏低，可能的原因是？

多选题 · hard

  • A.warp 正在等待全局内存数据← 正确答案
  • B.共享内存发生 bank 冲突← 正确答案
  • C.理论峰值本来就低
  • D.Tile 边缘存在无效计算浪费← 正确答案

答案：ABD。等待数据、bank 冲突、无效计算都会让锅虽热但没菜可炒或炒了白炒。理论峰值是固定属性，不会因为发射速率低而“本来就低”；反过来，峰值越高越说明有优化空间。

Q9. 在共享内存数组中加入 padding 的主要作用是什么？

单选题 · easy

  • A.减少 bank 冲突← 正确答案
  • B.提高 FP32 精度
  • C.增加 L2 缓存容量
  • D.降低寄存器占用

答案：A。共享内存被分成多个 bank，padding 可以把同一行的数据错开到不同 bank，避免厨师们抢同一个格子。它跟精度、L2 容量、寄存器占用都没有直接关系。

Q10. 在优化 Transformer 推理时，哪些做法能真正提升 Tensor Core 利用率？

多选题 · extreme

  • A.动态 batch 把请求拼成大矩阵← 正确答案
  • B.使用 FP16/INT8 混合精度← 正确答案
  • C.把 QK^T 矩阵乘交给 CPU 执行
  • D.在 kernel 中增加更多同步点

答案：AB。动态 batch 让细长条矩阵变成丰满方阵，FP16/INT8 让 Tensor Core 直接加速，都是提高利用率的有效手段。把 QK^T 给 CPU 会引入大量拷贝开销；增加同步点只会让锅频繁空烧，效果相反。