为什么。Triton >= 3.3 在 sm_80 以下会把 tl.dot 降级为标量 FMA（triton-lang/triton#5066）。因此在 V100 上，Triton 模式下每个矩阵乘法都跑在 CUDA core 上：引擎的 flash attention 在 T = 79 200、D = 96 时测得 2.46 TFLOP/s，约相当于一次 Allegro denoise 步骤耗时的 94%。我们改为从自己写的 Gluon kernel 访问 V100 的 tensor core（HMMA.884），且不在 Triton 之外引入原生代码、不依赖外部库、不额外加载 cubin、也不降级 Triton：Gluon（triton.experimental.gluon，Triton 3.8）没有 sm_70 限制，而 ptxas 接受 mma.sync.aligned.m8n8k4.row.col.f32.f16.f16.f32 用于 sm_70。

选择方式。硬件 profile 声明使用的单元；代码中不点名任何显卡：

src/neurobrix/config/vendors/nvidia/volta.yml

matrix_unit: shape: m8n8k4 operand_dtype: float16 accumulator_dtype: float32 flash:

• head_dim_le: 256 block_m: 64 block_n: 32 kernels/ops/_configs.py 读取它（matrix_unit()、matrix_unit_flash_tile(head_dim)）；SDPA wrapper 在 profile 声明了该 shape、一个 tile 行覆盖 head dim 且操作数处于声明的 dtype 时，采用 m8n8k4 flash。没有 matrix_unit 的 profile 则保持 tl.dot kernel 不变。

来自 CUTLASS cute/atom/mma_traits_sm70.hpp，SM70_8x8x4_F32F16F16F32_TN。一个 lane 是 a0 + 2 a1 + 4 q0 + 8 q1 + 16 b；四个 quadpair q 各自计算一个 16x16 tile 的一个 8x8 象限。某个 quadpair 的线程 t = a0 + 2 a1 + 4 b 持有 A 的第 t 行和 B 的第 t 列，对应 k = 0..3；其累加器寄存器 v = v0 + 2 v1 + 4 v2 持有 D[m = a0 + 2 v1 + 4 b, n = v0 + 2 a1 + 4 v2]。

一个 [BM, BN] 累加器按 [MB, 2, 8, NB, 2, 2, 4] = (mb, mi3, mi2..0, nb, ni3, ni2, ni1..0) 存放，其中 m = 16 mb + mi，n = 16 nb + ni。它的行主序就是 [BM, BN] 的行主序，所以 reshape 不移动任何数据。在这些坐标下，A fragment 只依赖 (mb, mi3, mi2..0, ni1..0)，B fragment 只依赖 (mi2..0, nb, ni3, ni1..0)：每个操作数只是 [MB, 2, 8, 1, 1, 1, 4] 和 [1, 1, 8, NB, 2, 1, 4] 的普通广播，由 inline_asm_elementwise 自行广播（pack = 8，每次调用八个累加器元素）。因此一个 fragment 寄存器可供给共享它的每个子 tile，无需 shuffle。

累加器布局是根据该 fragment 写出的 DistributedLinearLayout（mma_m8n8k4.acc_layout）：寄存器基为 ni0、mi1、ni2，随后是 NB 和 MB 的重复；lane 基为 mi0、ni1、mi3、ni3、mi2；warp 基先 MB 后 NB。

全局 tile 以合并方式加载，并以 [k-step x rows, 4] 暂存到共享内存（k 在最内层，每行 8 字节，按 LDS.64 读回）。一个 k-step 是二维 buffer 上的一个 .slice(row0, rows)，按从累加器布局派生的线性布局加载（frag_a_layout、frag_b_layout），然后在寄存器中 reshape 并 permute 成广播 shape；convert_layout(assert_trivial=True) 证明这种重新标记不移动数据。过程中满足的约束：共享 buffer 是二维的（rank-3 的 .index() 保持 rank-3 顺序），每个维度都是 2 的幂，全局常量用 gl.constexpr 包起来。

bank conflict。操作数的 k 沿行方向时（matmul 的 B、attention 的 V），用 BlockedLayout([4, 2], [1, 32], [nw, 1], [1, 0]) 加载，这样每个线程持有某一列中四个连续的 k，转置存储就是一次 STS.64。否则存储会有 16 路冲突：39 -> 58 TFLOP/s。

kernels/ops/flash_attention_m8n8k4.py 保持 flash_attention.py 的约定（Q 按 softmax_scale 缩放并舍入到其自身 dtype、常驻内存的加性 bias、GQA 分组、存储 log-sum-exp、尾部掩码、整行全掩码交给 wrapper 的 guard）。一个 warp 负责 16 个 query 行，因此行统计从不跨 warp；S = Q K^T 和 O += P V 是 HMMA；在线 softmax 在累加器布局上以 log2 域运行（行统计是其四个列维度归约）。head dim 被拆成至多两个 2 的幂（96 = 64 + 32），绝不补齐到下一个 2 的幂。

引擎的 launcher 绑定 kernel 时不用 JITFunction.create_binder，而 Triton 正是在那里为 @gluon.jit kernel 选择 GluonASTSource；launcher 做出同样的选择（launcher._source），否则模块缺少 ttg.num-warps，layout 验证会失败。

该时钟下的 HMMA 上限：80 SM x 8 tensor core x 128 FLOP/clk x 1.290 GHz = 105.7 TFLOP/s（125 这个数字是 1530 MHz boost 下的值）。

产生这些结果的探针在 tools/tensor_cores_m8n8k4/ 中，连同其扫描输出。