这些数字是怎么来的¶
每个数据页只回答一个问题：在同一个工作负载上，TileOPs 与同一算子最快的其它实现相比如何？ 每个算子一张表，每个工作负载一行。不做跨工作负载的平均 —— 页面上每一个数字都只属于一个具体的 shape 和 dtype。
对照的是谁¶
比值 = 基准耗时 ÷ 我们的耗时，在同一测量口径下，按每个 (算子, 工作负载, dtype) 分别计算。两边都按该行的 shape 和 dtype 选择实现。
基线文档决定哪些实现有资格进入对照组。实测在确切工作负载上，选择主测量口径下符合准入条件的最快候选。「对照实现」列保留各候选，所选基准标为 基准。
「对照实现」列保留每个候选的名称、来源档位和实测耗时，包括开源库候选。每行只显示主测量口径的一个比值。标为 eager 的 MLIR 候选未经 graph 捕获，其耗时不替代 graph 主比值的基准。
只有一个候选的池，harness 记作 single_candidate：那个工作负载上只建得起一个对手。它仍然是一次实测比较 —— 只是没有可比的第二家。
档位是什么意思¶
对照实现 里的每一行都带一个档位。档位记录的是这个 kernel 的来源，它完全不说明谁更快。
| 档位 | 含义 |
|---|---|
| 开源库 | 从第三方 Ascend 算子库的源码编出来的 kernel，通过那个库自己的入口点调用。只有拿到「该库自己编出来的 kernel 确实跑了」的证据才被接纳：一个缺 kernel 的自定义算子包会静默回落到 CANN 内置，调用照样成功、输出照样正确、时间照样看着合理。所以这里的来源认定靠的是追踪进程实际打开了哪个二进制，绝不是「调用没报错」。 |
| tilelang-ascend | tilelang-ascend（Tile-AI）原始示例或测试 kernel，与我们使用相同 DSL/编译器；通过源码来源证据和数值门禁后才进入实测候选池。此档位说明 kernel 来源，不代表速度排名。 |
| 厂商库 | 厂商实现：CANN 内置算子，或 torch_npu 自己对这个算子的分发。它是完全相同工作负载上的一个真实实现，而且在这块硬件上经常就是最快的那个。 |
| torch_compile | 把同一段 PyTorch 程序交给这块后端上的 torch.compile（aclgraph 或 ge）编出来的实现，作为候选池里的一员一起实测。它既不是手写 kernel，也不是厂商分发，而是编译器对同一个问题给出的答案。正因为两边都不属于，它用的是中性徽章；而它在这里赢的次数不少，读者有必要知道赢我们的到底是什么。 |
| Agent 编写 kernel | 由 agent 编写的 kernel，按 D060 与其它档位同样的源码来源证据和数值门禁进入候选池。和这张表里每一个档位一样，它说明的是 kernel 的来源，不是它的速度。当前快照中没有任何一行属于这个档位。 |
所以同一行里 开源库 的时间比 厂商库 慢，不是错误；在这里的好几个算子族上，这就是常态。把这个徽章读成强弱排名，正是这一列存在的目的所要防止的那个误解。
历史普查的分母为 91 个 elementwise、reduction、scan、dropout 算子。按本页所渲染的这份快照当场算出：32/91 个至少有一个 handwritten 或 tilelang_ref 来源候选取得实测时间，58/91 个没有，另有 1/91 个因本次运行没有发布任何 workload 而无法判定。这是 shape/dtype 契约门之后的实测候选池覆盖，不是穷尽源码普查，也不是数值正确性通过率。逐 workload 的实际候选与拒绝原因以记录为准；只赢过厂商实现不能证明赢过开源 kernel。
这一行是在哪个口径下测的¶
一个比值只和同口径的另一个比值可比。这里每个工作负载都在一个具名的测量口径（regime）下测得，本次运行的主口径是 graph —— 覆盖情况那几条、颜色、算子排序，说的全是这个口径。
有几个算子在这个口径下压根没有测量，发布出来的是一份只测过 eager 的旧记录。这些行保留 —— 那是一次真实工作负载上的真实测量，删掉会让页面凭空变短而且不给理由。取而代之的做法是把它标出来。
eager — 这一行的比值是在 eager 下测的，不是 graph。它只能和其它 eager 行比，绝不能和它上下的那些行比。
概览页上的每一个计数都守同一条规矩：graph 口径下的算子数和其它口径下的算子数分成两个数报，绝不合并成一个。
颜色就是结论¶
| 含义 | |
|---|---|
| 0.74× | 比对照实现慢 —— 低于 0.95×。 |
| 1.02× | 与它持平 —— 0.95–1.05×，落在测量噪声内。 |
| 1.42× | 比它快 —— 1.05× 及以上。 |
| 18.06× | 没有快的对照实现可比 —— 只有一个功能参考实现（名字以 -ref 结尾）。这个数字意义不大。 |
| — | 这个工作负载上没有任何对照实现跑过。 |
比值 = 对照实现的耗时 ÷ 我们的耗时，所以大于 1 表示 TileOPs 更快。
各列含义¶
| 列 | 含义 |
|---|---|
| 工作负载 | W1、W2、… —— 每张表上方的图例会把每一个展开：benchmark 自己给它的 id、它跑的 dtype，以及每个输入张量（写成 名称: shape, dtype）。形状相同的张量并列在一起，但各自带自己的 dtype，所以一个 bool 的 mask 会在被读到的地方就标明。张量之后是那些决定算子规模但不决定形状的维度（GEMM 的 m/n/k，MoE 路由的 num_experts），再往后是灰色的、调用时没有沿用签名默认值的参数。已经能由其它量确定的不再重复 —— 例如 max_seqlen_q 就是 max(q_lens)。 |
| 比值 | 对照 / 我们 —— 基准的耗时除以我们的耗时。颜色评的就是这一个数。 各候选的名称、档位和逐候选耗时见「对照实现」列。 |
| 耗时 | 本次调用在 device 上执行它各个 kernel 的区间并集，单位毫秒。本页所有比较都用它。每个工作负载另存了一份 host 挂钟读数作端到端参考 —— 两者之差约 40–50 µs，所以便宜的工作负载用 host 计时会把比值推向 1.0。详见「测量方法」一节。 |
| 对照实现 | 每个候选占一行，各自带耗时（ms）和来源档位。标为 基准 的候选用于主比值。名称缩写可悬停查看完整绑定。不同测量口径的耗时应按各自口径解读，包括标为 eager 的候选。若运行未发布候选池，则显示实测的具名基线：调优过的库 kernel（fla、mamba、fa3、triton 等）、PyTorch 原生算子（torch），或名字以 -ref 结尾的 eager 参考拼装实现。 |
| 吞吐 | TFLOP/s：所需 FLOPs ÷ 耗时。这个 FLOP 数是解析算出来的 —— 用算子的 eval_roofline 公式代入该工作负载自己的 shape，不是硬件计数器 —— 所以它算的是问题本身要求的工作量，不是 kernel 实际发出的指令。padding、重算、被 mask 掉的 tile 在这里都看不见；这个数只在同一算子、同一工作负载的不同实现之间可比。 |
| SOL | 占算法光速（speed-of-light）的比例：该工作负载在物理上最快可能的时间 ÷ 我们的耗时。比值 那一列说的是今天有没有人比我们快；SOL 说的是任何人最多还能快多少。详见下文。 |
| 瓶颈 | 决定这个工作负载下限的资源：mem（HBM 搬运）、comp（计算吞吐），或 lat —— 工作负载太小，模型无法判定，它的 SOL 数字也会随之置灰。 |
这个耗时具体怎么测的 —— 算进了什么、漏掉了什么、什么情况下它拒绝给出数字 —— 见 Benchmark 怎么计时。
每个算子的标题上带着它的工作负载数和测试结果（✅ 通过 · ❌ 失败 · ⏭️ 全部跳过 · · 没有匹配到测试）。
光速（Speed of light）¶
SOL 是算法意义上的光速效率：max(字节数 / 带宽, FLOPs / 计算屋顶) / 耗时，其中的天花板用的是这台机器标定过的值 —— 微基准在这块硬件上实际达到的带宽和计算速率，不是规格书上的数。100% 意味着这个算法在这块硬件上不可能有更快的实现。
三句话界定一个 SOL 读数的含义：
- 字节数是算法的最小搬运量 —— 每个输入读一次、每个输出写一次 —— 不是 kernel 实际产生的 DRAM 流量。一个把数据搬两遍的 kernel 会得低分；这正是它该得低分的地方。
- FLOPs 按 TileOPs 的计数约定（一个超越函数算一次），不是按指令的硬件代价。所以这个指标不会把一个受特殊函数单元限制的 kernel 认证为已到极限。
- 计算屋顶取的是「最优实现会用哪个单元」 —— 每个算子显式声明，绝不从正在跑的 kernel 反推。所以一个跑在错误单元上的 kernel，衡量它的仍然是那个正确的天花板。
| SOL | 瓶颈 | 含义 |
|---|---|---|
| 92% | mem | 已到可达天花板（自 80% 起）。这条天花板是对各种访存组合取的包络，而一个 kernel 自己的组合会低于包络，所以这条线给每种组合都留了余量。再优化最多也只能拿到剩下那一点。 |
| 63% | mem | 还有余量。 |
| 41% | lat | 工作负载太小，模型无法判定 —— 主导测量的是启动开销，不是 roofline —— 所以数字照给，但不评级。 |
| ⚠ 108% | mem | 高于标定的天花板：说明公式或标定有一个是错的。绝不能读成「这是个快 kernel」。 |
· |
· |
缺少某个输入：没有 roofline 公式、计时方式不是 device 侧采集，或者这个设备没有硬件档案。 |
这个模型、它的阈值以及公式审计机制，规定在 TileOPs 的 docs/design/roofline.md 里；本页直接导入那份实现，而不是自己重新推导一遍。
shape 是从哪来的¶
shape 是从 TileOPs 的 spec manifest 里读出来的，按 benchmark id 里的 label 和 dtype 关联到对应行 —— 所以一行给出的是这个算子声明可以接受的形状。manifest 没有声明的工作负载 —— 也就是人工编写的、不由 spec 驱动的边界用例探针 —— 改为显示快照为那次运行记录下来的形状：那是实测的一次调用，不是一份声明。两者都没有时，这一行只显示 benchmark id，下面没有 shape。
空单元格¶
· 表示这个指标的某个输入没有被记录，绝不表示值是零：可能是该算子在这个工作负载上没有报告 FLOP 数，或者根本没有对照实现在它上面跑过。
在对照组里，同一种「缺失」有它自己的写法。被选中却从未被实测的候选，记作 not-timed。如果本次运行在别处发布了那个基线的时间，单元格就显示那个值并跟一个 *，把原因写在单元格的提示里；如果连别处也没有，单元格就是 ·。两种情况都绝不会渲染成 0 —— 那会被读成一个无限快的 kernel。