第12章 12Fusion：融合算子

本章导读

前两章已经能分别完成矩阵乘和 Softmax。把它们接起来后，前一步的输出就成了后一步的输入：先算完、写入数组，再由下一步读走，是最直接的实现。

如果后一步只需要中间结果的一部分，能否在这部分还可用时继续计算，省去完整中间数组？本章用 Attention 研究这个问题。先用三个 key 手算一行输出，再让已经算过的贡献随着新 key 到来继续累积，最后用完整路径的计时判断这项改动是否值得。

配套代码 编程范式附录

当前环境Radeon RX 9070 XT · ROCm 10.0 · 原生 Ubuntu 24.04

HIP 深入路线从线程分工到访存与协作Triton 快速路线用数据块表达同一个算子

• 公共基础：阅读加权和、中间数据和在线更新，看清结果为什么不变。
• HIP 路线：从三个 kernel 的物化基线出发，比较一个 block 完成一行的在线版本，检查减少中间读写是否抵得上新增协作。
• Triton 路线：从三个 kernel 的物化基线出发，先比较在线计算，再固定其他参数，只改变一次处理的 key 数。

配套代码在 code/part2-kernels/chapter12/。环境沿用 Radeon RX 9070 XT（gfx1201）+ ROCm 10.0 + 原生 Ubuntu 24.04。动画展示算法中的依赖和手算过程，不表示 GPU 的真实执行速度。

12.1 Attention 的一行加权和

本章计算单组 Q、K、V 的 FP32 前向 Attention，不包含 batch、多 head 或 causal mask。序列长度为 $S$，每个向量有 $D$ 个分量，三个输入的形状都是 [S,D]。

固定第 $i$ 行 query，先与每行 key 做点积得到分数，再对这一行分数做 Softmax，最后用这些权重合并 value：

$$s_{ij}=\frac{Q_i\cdot K_j}{\sqrt{D}},p_{ij}=\frac{e^{s_{ij}}}{\sum _{k=0}^{S-1}{e}^{s_{ik}}},O_i=\sum _{j=0}^{S-1}{p}_{ij}{V}_j.$$

一个 $p_{ij}$ 是标量，$V_j$ 是向量；它们相乘后仍有 $D$ 个分量。把所有 key 对应的贡献加起来，得到这一行输出 $O_i$。因此，计算某个 query 行需要遍历全部 key/value，但不需要其他 query 行的结果。

先假设某个 query 得到三个分数 [2,1,4]，对应的 value 是 [1,0]、[0,2]、[3,1]。沿用第 10 章的稳定 Softmax，先减去最大值 4：

key	平移后的指数，约值	归一化权重，约值	对输出的贡献，约值
0	$e^{-2}=0.1353$	0.1142	[0.1142, 0]
1	$e^{-3}=0.0498$	0.0420	[0, 0.0840]
2	$1$	0.8438	[2.5314, 0.8438]

未舍入指数的总和约为 1.1851，输出约为 [2.6456,0.9278]。后面始终跟踪同一组数据，观察怎样在不保存全部分数的情况下得到这个结果。这个例子用于手算，不是性能实验输入。

先在 图 12.1 中逐项观察：一个概率分别乘对应 V 的两个分量，得到两个贡献，再各自加入输出的累计值。概率和 V 的原值始终保留，只有三个 key 都处理完后才写出输出。

ATTENTION · 一行手算一个概率怎样贡献两个输出分量

1/6

固定一个 query 的三个 key三个 score 是 2、1、4；每个 V 都有两个分量。先保留这些输入，输出还未写出。

将三个 score 转成概率减去最大值 4，计算指数再除以指数和。得到三个概率，它们各自用于对应 V 的两个分量。

第一项贡献进入两个累计值p₀ 乘 V₀=[1,0]，得到 [0.1142,0]。两个分量分别累加；源概率和 V 不会被移走。

第二项只增加分量 1p₁ 乘 V₁=[0,2]，得到 [0,0.0840]。零分量也按同一乘法规则处理。

第三项贡献加入已有累计p₂ 乘 V₂=[3,1]，得到 [2.5314,0.8438]，分别加到两个已有累计值上。

三个 key 全部处理完，写出 O最终 O≈[2.6456,0.9278]。一个概率是标量，一次乘法却贡献整个向量。

固定一个 query 的三个 key三个 score 是 2、1、4；每个 V 都有两个分量。先保留这些输入，输出还未写出。

步骤 1 / 6输入

数学步骤示意；数字不插值，动画速度不表示 GPU 耗时。

图 12.1　同一个 query 的三个权重分别作用于三条 V。可在每个 key 的步骤停留，核对两个乘积和两个累计分量；动画的播放时长不表示计算耗时。

12.2 中间矩阵与物化成本

对所有 query 行重复上述计算，就会得到分数矩阵 Scores[S,S] 和概率矩阵 P[S,S]。把一个中间结果真正保存成数组，称为物化（materialize）：

Q、K → Scores [S,S] → P [S,S] → 与 V 合并 → O [S,D]

两条路线都先建立三个 kernel 的物化基线。分数、Softmax、加权和各自完成一段工作，通过全局数组传递结果。每段职责清楚，也便于定位错误。

① 分数

一个 Q 行与全部 K逐行配对

写出 Scores 行[2, 1, 4]

→

② Softmax

读取完整 Scores 行

写出 P 行p0=0.1142p1=0.0420p2=0.8438

→

③ 加权和

完整 P 行与全部 V逐项合并

写出 O 行O[0]=2.6456O[1]=0.9278

每个 query 都重复这条依赖：Scores / P 的完整形状为 S×S，O 为 S×D。箭头表示数组交接，不表示 GPU 调度或时间。

图 12.2　三个阶段通过 Scores 与 P 交接。仍固定上例的一个 query：分数需要全部 K，输出需要完整概率行和全部 V。每个 query 都按此规则得到自己的输出行；箭头只表示数组依赖。

图 12.2 中，两块 FP32 中间数组合计占用 $2S^2\times 4$ Byte。对主形状 S=128,D=64，它们的容量合计为 128 KiB。只按每块数组完整写一次、完整读一次计，已有 256 KiB 的逻辑访问。这是按数组大小计算的最低账目，不是测得的带宽或流量。

具体源码可能重复读取：例如 HIP 的加权和 kernel 由一个线程负责一个输出分量，同一行概率会被该行的多个输出线程分别读取；行 Softmax 也会多次扫描 Scores。缓存还会影响这些访问最终到达哪一级存储。因此，不能拿 16S² Byte 直接代替本程序的全部读取量，再用总时间反推显存带宽。

需要检验的假设是：省去完整 Scores/P 的中间读写，能否降低一次 Attention 的总时间？ 为此必须改变计算过程，让尚未读完 key 时也能保留正确的历史贡献。

12.3 在线 Softmax 与历史贡献

先只处理第一个分数 2。以它作为当前最大值，指数就是 1，加权 value 是 [1,0]。第二个分数 1 小于当前最大值，加入 $e^{-1}$ 和 $e^{-1}[0,2]$ 即可。

第三个分数变成 4。之前两个指数以 2 为基准，现在必须一起换成以 4 为基准：旧结果都乘 $e^{2-4}$，再加入新的 [3,1]。不能只重缩放分母而保留原来的分子。

ATTENTION · 一行手算新最大值到来后，保留并更新历史贡献

1/7

为同一行保留三个状态m 是最大值，l 是指数和，a 有两个加权累计分量。初始为 −∞、0、[0,0]。

以第一个 score 为基准score=2，指数为 1。得到 m=2、l=1、a=[1,0]；此时还不能当作最终输出。

最大值不变，加入较小项score=1，指数为 exp(−1)。得到 l≈1.3679、a≈[1,0.7358]，仍保留基准 m=2。

新 score=4，先保留已有贡献旧 l 和两个旧 a 都基于 m=2。新最大值变为 4，下一步必须把所有历史项一起换基准。

旧 l 与两个旧 a 同乘 exp(−2)得到 l≈0.1851、a≈[0.1353,0.0996]。旧值仍在左侧可核对，此步尚未加入 key 2。

新项指数为 1，加入 V₂在重缩放后的 l 上加 1，两个 a 分别加 3 和 1。得到 l≈1.1851、a≈[3.1353,1.0996]。

两个分量分别除以同一个 l所有 key 处理完后，O=a/l≈[2.6456,0.9278]，与直接概率加权和一致。

为同一行保留三个状态m 是最大值，l 是指数和，a 有两个加权累计分量。初始为 −∞、0、[0,0]。

步骤 1 / 7初始化

数学步骤示意；数字不插值，动画速度不表示 GPU 耗时。

图 12.3　固定同一个 query，依次加入三个 key 的贡献。新最大值从 2 变为 4 时，分别停在“保留旧值”“重缩放”“加入新项”三个步骤，核对旧分母与两个旧累计是否同乘 exp(−2)。最后再做除法；显示的小数只用于阅读，后续计算使用未舍入的值。

把 图 12.3 中持续保留的值记为三个状态：

状态	当前含义	需要保存多少数据
$m$	已处理分数的最大值	一个标量
$l$	以 $m$ 为基准的指数和	一个标量
$a$	以同一基准累计的加权 value	一个 $D$ 维向量

初始状态为 $m=-\mathrm{\infty },l=0,a=0$。读到分数 $s_j$ 后，令

$$m^{\prime }=max(m,s_j),\alpha =e^{m-m^{\prime }},\beta =e^{s_j-m^{\prime }}.$$

新的分母和分子分别为

$$l^{\prime }=\alpha l+\beta ,a^{\prime }=\alpha a+\beta V_j.$$

最后输出 $a/l$。重缩放来自恒等式 $e^{s-m^{\prime }}=e^{s-m}{e}^{m-m^{\prime }}$，它让所有历史项仍采用同一尺度。第一个 key 有限时，$\alpha =e^{-\mathrm{\infty }}=0$，初始状态不会贡献任何权重。

这称为在线（online）计算：随着新数据到来更新已有状态。也可以一次处理一个 key tile，先求 tile 内的最大值、指数和与加权贡献，再按同样的关系合入历史。两种组织在数学上等价；FP32 的归约顺序不同，最后几位不要求逐位相同。

12.4 正确性参考与计时范围

两条路线都使用连续的 FP32 Q/K/V[S,D]，输出为 FP32 O[S,D]。性能实验输入按行号、分量号和固定 seed 生成；所有值都是整数除以二的幂，因此 HIP 与 Triton 得到相同的输入，不依赖两种随机库恰好生成相同序列。

检查项目	本章约定
主形状与输入	S=128,D=64，input=dyadic，seed=20260920
正确性参考	CPU 用 FP64 完成 QK、稳定 Softmax、PV，最后转为 FP32
逐元素误差	$
有限值	任意输出为 NaN 或无穷立即失败
检查时机	第一次输出预填 NaN，计时前检查完整输出；计时后再检查最后一次完整输出
GPU 算术	两路线保持 FP32；本章 Triton 不使用 tl.dot 或混合精度矩阵路径

六个配置分别检查 dyadic 输入的 1×1、17×31、33×65、65×64，以及 rising-max 输入的 65×64、129×64，共 36 项检查通过。形状顺序为 S×D；其中 33/65/129 检查 key 尾部，31/65 检查向量分量尾部。

专用输入 rising-max 让后一个 key 的分数高于前一个，迫使状态在多个 key、多个 tile 之间反复重缩放。它用于发现“只有最大值不变时才算对”的错误；不与默认输入的性能样本混在一起。单 key 检查还有一个直接答案：权重为 1，输出应等于这一行 value。

GPU event 计完整前向路径：物化版本包含三个 kernel，在线版本包含一个 kernel。分配、输入拷贝、CPU 参考、正确性检查及 Triton 首次编译在稳态计时之外。正式计时不启用 profiler，诊断另行采集。

每配置运行 3 个独立进程，每进程预热 10 次、采样 50 次，每个 event 区间后同步。图的中心值是三个进程中位数的中位数，范围线是它们的最小—最大值，不是置信区间。主形状与两组附加形状共 54 个进程、2700 个样本；扫描后另外运行 18 个进程、900 个样本作独立确认。

两个 Host 当前都为 Scores/P 分配空间，即使只运行在线版本也保留这项分配。输出中的 intermediate_allocated_bytes 记录这两块缓冲区容量，intermediate_used_bytes 则在在线版本中为 0。后一个字段表示 kernel 不使用它们，既不是访存计数器，也不能证明整个进程的显存分配峰值已下降。

12.5 HIP 与 Triton 优化路线

先在一种语言内建立物化基线，再对照在线实现。两条路线分别回答自己的优化问题，不把彼此当作必须战胜的初版。

环境与结果目录。 以下命令在 GPU 机器的仓库根目录开始执行，后续保持在 code/part2-kernels/。沿用本篇 uv.lock，不要删除锁文件：

cd code/part2-kernels &&
uv sync --locked &&
source ./activate-rocm.sh

为本章新建一个独立结果目录，保存打印出来的路径：

mkdir -p chapter12/results &&
RUN=$(mktemp -d chapter12/results/walkthrough-XXXXXX) &&
mkdir "$RUN/build" &&
printf '本轮结果目录：%s\n' "$RUN"

换终端后，重新执行环境准备，再恢复之前打印的目录；提示处只粘贴路径，不含 RUN=：

read -r -p '粘贴上次输出的结果目录：' RUN &&
test -n "$RUN" && test -d "$RUN/build" &&
printf '继续使用结果目录：%s\n' "$RUN"

需要重测或改动源码时，新建 RUN，重新编译 HIP；恢复旧目录只用于继续未修改的实验。后续命令会拒绝写入已有样本。

每条命令先校验，再完成一个进程的 10 次预热与 50 次计时；折叠输出对应第 1 个进程。图使用三个独立进程的统计，不是一次终端输出。 手工复测时同时把样本名 p1 与 --process 1 改为 p2/2、p3/3，每遍覆盖同组候选并轮换起始项。章末批量入口会自动完成这一步、汇总并画本轮图，不要求先运行 profiler。

先判断完整结果与完整时间。HIP 在线版本出现明显退步，才进一步追问设备上的哪个阶段变长；Triton 的采用与 tile 选择先由重复测量判断，不要求两条路线走相同的工具流程。

切换会保留各自的阅读位置

12.5.1 HIP 基线：分阶段保存中间结果

完整源码：code/part2-kernels/chapter12/attention_hip.hip

#include <hip/hip_runtime.h>

#include <algorithm>
#include <cmath>
#include <cstdlib>
#include <cstdint>
#include <fstream>
#include <iomanip>
#include <iostream>
#include <limits>
#include <numeric>
#include <string>
#include <vector>

#define HIP_CHECK(call)                                                        \
    do {                                                                       \
        hipError_t error = (call);                                             \
        if (error != hipSuccess) {                                             \
            std::cerr << "HIP error: " << hipGetErrorString(error) << '\n';   \
            std::exit(1);                                                      \
        }                                                                      \
    } while (0)

__global__ void scores_kernel(const float* q, const float* k, float* scores,
                              int seq, int dim, float scale) {
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index >= seq * seq) return;
    int row = index / seq;
    int col = index % seq;
    float value = 0.0f;
    for (int d = 0; d < dim; ++d) {
        value += q[row * dim + d] * k[col * dim + d];
    }
    scores[index] = value * scale;
}

__global__ void softmax_rows_kernel(const float* scores, float* probabilities,
                                    int seq) {
    extern __shared__ float shared[];
    int row = blockIdx.x;
    int tid = threadIdx.x;
    float local_max = -INFINITY;
    for (int col = tid; col < seq; col += blockDim.x) {
        local_max = fmaxf(local_max, scores[row * seq + col]);
    }
    shared[tid] = local_max;
    __syncthreads();
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) shared[tid] = fmaxf(shared[tid], shared[tid + stride]);
        __syncthreads();
    }
    float row_max = shared[0];
    // Finish every read of the maximum before reusing LDS for local sums.
    __syncthreads();
    float local_sum = 0.0f;
    for (int col = tid; col < seq; col += blockDim.x) {
        local_sum += expf(scores[row * seq + col] - row_max);
    }
    shared[tid] = local_sum;
    __syncthreads();
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) shared[tid] += shared[tid + stride];
        __syncthreads();
    }
    float denominator = shared[0];
    for (int col = tid; col < seq; col += blockDim.x) {
        probabilities[row * seq + col] =
            expf(scores[row * seq + col] - row_max) / denominator;
    }
}

__global__ void probability_value_kernel(const float* probabilities,
                                         const float* value, float* output,
                                         int seq, int dim) {
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index >= seq * dim) return;
    int row = index / dim;
    int d = index % dim;
    float result = 0.0f;
    for (int col = 0; col < seq; ++col) {
        result += probabilities[row * seq + col] * value[col * dim + d];
    }
    output[index] = result;
}

__global__ void online_attention_kernel(const float* q, const float* k,
                                        const float* value, float* output,
                                        int seq, int dim, float scale) {
    extern __shared__ float shared[];
    float* reduction = shared;
    float* numerator = reduction + blockDim.x;
    float* state = numerator + dim;  // m, l, alpha, beta
    int row = blockIdx.x;
    int tid = threadIdx.x;

    for (int d = tid; d < dim; d += blockDim.x) numerator[d] = 0.0f;
    if (tid == 0) {
        state[0] = -INFINITY;
        state[1] = 0.0f;
    }
    __syncthreads();

    for (int col = 0; col < seq; ++col) {
        float partial = 0.0f;
        for (int d = tid; d < dim; d += blockDim.x) {
            partial += q[row * dim + d] * k[col * dim + d];
        }
        reduction[tid] = partial;
        __syncthreads();
        for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
            if (tid < stride) reduction[tid] += reduction[tid + stride];
            __syncthreads();
        }
        if (tid == 0) {
            float score = reduction[0] * scale;
            float old_max = state[0];
            float new_max = fmaxf(old_max, score);
            float alpha = expf(old_max - new_max);
            float beta = expf(score - new_max);
            state[0] = new_max;
            state[1] = state[1] * alpha + beta;
            state[2] = alpha;
            state[3] = beta;
        }
        __syncthreads();
        for (int d = tid; d < dim; d += blockDim.x) {
            numerator[d] = numerator[d] * state[2] +
                           state[3] * value[col * dim + d];
        }
        __syncthreads();
    }
    for (int d = tid; d < dim; d += blockDim.x) {
        output[row * dim + d] = numerator[d] / state[1];
    }
}

// Both runtimes use this element-wise contract against a CPU FP64 reference
// rounded once to FP32. NaN/Inf always fails, independently of the error value.
constexpr double kAtol = 2.0e-5;
constexpr double kRtol = 2.0e-4;

struct Options {
    std::string version = "all";
    std::string input = "dyadic";
    std::string samples;
    std::string config_id;
    int seq = 128;
    int dim = 64;
    int block = 256;
    int warmup = 10;
    int repeat = 50;
    std::uint64_t seed = 20260920;
    int process = 1;
};

[[noreturn]] void usage_error(const std::string& message) {
    std::cerr << message << '\n';
    std::exit(2);
}

std::uint64_t unsigned_integer(const std::string& text) {
    if (text.empty() || text.find_first_not_of("0123456789") != std::string::npos)
        usage_error("expected a non-negative integer: " + text);
    try {
        return std::stoull(text);
    } catch (const std::exception&) {
        usage_error("integer out of range: " + text);
    }
}

Options parse_options(int argc, char** argv) {
    Options options;
    for (int i = 1; i < argc; ++i) {
        const std::string name = argv[i];
        if (name == "--help") {
            std::cout << "--version all|materialized|online --seq S --dim D "
                         "--block THREADS --input dyadic|rising-max --seed UINT32 "
                         "--warmup N --repeat N --config-id ID --process N "
                         "--samples FILE\n";
            std::exit(0);
        }
        if (i + 1 >= argc) usage_error("missing value for " + name);
        const std::string value = argv[++i];
        if (name == "--version") options.version = value;
        else if (name == "--input") options.input = value;
        else if (name == "--samples") options.samples = value;
        else if (name == "--config-id") options.config_id = value;
        else if (name == "--seed") options.seed = unsigned_integer(value);
        else {
            const auto number = unsigned_integer(value);
            if (number > static_cast<std::uint64_t>(std::numeric_limits<int>::max()))
                usage_error("integer too large for " + name);
            const int integer = static_cast<int>(number);
            if (name == "--seq") options.seq = integer;
            else if (name == "--dim") options.dim = integer;
            else if (name == "--block") options.block = integer;
            else if (name == "--warmup") options.warmup = integer;
            else if (name == "--repeat") options.repeat = integer;
            else if (name == "--process") options.process = integer;
            else usage_error("unknown option: " + name);
        }
    }
    const bool valid_version = options.version == "all" ||
                              options.version == "materialized" ||
                              options.version == "online";
    const bool power_of_two = options.block > 0 &&
                             (options.block & (options.block - 1)) == 0;
    if (!valid_version || options.seq <= 0 || options.seq > 4096 ||
        options.dim <= 0 || options.dim > 256 || !power_of_two ||
        options.block > 1024 || options.repeat <= 0 || options.process <= 0 ||
        options.seed > 0xFFFFFFFFULL ||
        (options.input != "dyadic" && options.input != "rising-max"))
        usage_error("invalid arguments (1 <= seq <= 4096, 1 <= dim <= 256)");
    if (options.config_id.find_first_not_of(
            "abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789-_.") != std::string::npos)
        usage_error("invalid config-id");
    if (!options.samples.empty() &&
        (options.version == "all" || options.config_id.empty()))
        usage_error("--samples requires one version and a --config-id");
    return options;
}

struct Inputs {
    std::vector<float> q;
    std::vector<float> k;
    std::vector<float> value;
};

Inputs make_input(const Options& options) {
    const std::size_t size = static_cast<std::size_t>(options.seq) * options.dim;
    Inputs input{std::vector<float>(size), std::vector<float>(size),
                 std::vector<float>(size)};
    for (int row = 0; row < options.seq; ++row) {
        for (int d = 0; d < options.dim; ++d) {
            const std::size_t index = static_cast<std::size_t>(row) * options.dim + d;
            const auto r = static_cast<std::uint64_t>(row);
            const auto column = static_cast<std::uint64_t>(d);
            input.q[index] = (static_cast<int>((17 * r + 13 * column + options.seed) % 65) - 32) / 64.0f;
            input.k[index] = (static_cast<int>((29 * r + 7 * column + options.seed + 11) % 65) - 32) / 64.0f;
            input.value[index] = (static_cast<int>((11 * r + 19 * column + options.seed + 23) % 65) - 32) / 32.0f;
            if (options.input == "rising-max") {
                // Every later key has a larger score, including across tiles.
                input.q[index] = d == 0 ? 8.0f : 0.0f;
                input.k[index] = d == 0 ? (row - options.seq / 2) / 8.0f : 0.0f;
            }
        }
    }
    return input;
}

std::vector<float> reference_attention(const Inputs& input, int seq, int dim) {
    std::vector<float> output(static_cast<std::size_t>(seq) * dim);
    std::vector<double> weights(seq);
    const double scale = 1.0 / std::sqrt(static_cast<double>(dim));
    for (int row = 0; row < seq; ++row) {
        double maximum = -std::numeric_limits<double>::infinity();
        for (int col = 0; col < seq; ++col) {
            double score = 0.0;
            for (int d = 0; d < dim; ++d)
                score += static_cast<double>(input.q[row * dim + d]) * input.k[col * dim + d];
            weights[col] = score * scale;
            maximum = std::max(maximum, weights[col]);
        }
        double denominator = 0.0;
        for (double& weight : weights) {
            weight = std::exp(weight - maximum);
            denominator += weight;
        }
        for (int d = 0; d < dim; ++d) {
            double result = 0.0;
            for (int col = 0; col < seq; ++col)
                result += (weights[col] / denominator) * input.value[col * dim + d];
            output[row * dim + d] = static_cast<float>(result);
        }
    }
    return output;
}

struct Validation {
    bool correct = true;
    bool finite = true;
    double max_absolute_error = 0.0;
    double max_tolerance_ratio = 0.0;
};

Validation validate(const std::vector<float>& actual, const std::vector<float>& reference) {
    Validation result;
    for (std::size_t i = 0; i < actual.size(); ++i) {
        if (!std::isfinite(actual[i]) || !std::isfinite(reference[i])) {
            result.correct = false;
            result.finite = false;
            result.max_absolute_error = std::numeric_limits<double>::infinity();
            result.max_tolerance_ratio = std::numeric_limits<double>::infinity();
            continue;
        }
        const double error = std::abs(static_cast<double>(actual[i]) - reference[i]);
        const double tolerance = kAtol + kRtol * std::abs(static_cast<double>(reference[i]));
        result.max_absolute_error = std::max(result.max_absolute_error, error);
        result.max_tolerance_ratio = std::max(result.max_tolerance_ratio, error / tolerance);
        result.correct = result.correct && error <= tolerance;
    }
    return result;
}

struct Timing {
    double minimum_ms;
    double median_ms;
    double mean_ms;
    std::vector<double> samples;
};

template <typename Launch>
Timing benchmark(const Options& options, Launch launch) {
    for (int i = 0; i < options.warmup; ++i) launch();
    HIP_CHECK(hipDeviceSynchronize());
    hipEvent_t start, stop;
    HIP_CHECK(hipEventCreate(&start));
    HIP_CHECK(hipEventCreate(&stop));
    std::vector<double> samples;
    samples.reserve(options.repeat);
    for (int i = 0; i < options.repeat; ++i) {
        HIP_CHECK(hipEventRecord(start, nullptr));
        launch();
        HIP_CHECK(hipEventRecord(stop, nullptr));
        HIP_CHECK(hipEventSynchronize(stop));
        float elapsed = 0.0f;
        HIP_CHECK(hipEventElapsedTime(&elapsed, start, stop));
        samples.push_back(elapsed);
    }
    HIP_CHECK(hipEventDestroy(start));
    HIP_CHECK(hipEventDestroy(stop));
    auto sorted = samples;
    std::sort(sorted.begin(), sorted.end());
    const std::size_t middle = sorted.size() / 2;
    const double median = sorted.size() % 2 == 0
        ? (sorted[middle - 1] + sorted[middle]) / 2.0 : sorted[middle];
    return {sorted.front(), median,
            std::accumulate(samples.begin(), samples.end(), 0.0) / samples.size(), samples};
}

void print_result(const Options& options, const std::string& version,
                  const Validation& validation, const Timing* timing,
                  const char* precheck, const char* postcheck) {
    const std::size_t intermediate_bytes = 2ULL * options.seq * options.seq * sizeof(float);
    std::cout << std::setprecision(17)
              << "RESULT operator=attention implementation=hip-" << version
              << " runtime=hip config_id=" << options.config_id
              << " shape=" << options.seq << 'x' << options.dim
              << " seq=" << options.seq << " dim=" << options.dim
              << " dtype=float32 input_precision=fp32 block=" << options.block
              << " block_k=" << (version == "online" ? 1 : 0)
              << " block_d=0 num_warps=0 launches=" << (version == "materialized" ? 3 : 1)
              << " warmup=" << options.warmup << " repeat=" << options.repeat
              << " seed=" << options.seed << " process=" << options.process
              << " timing=gpu-event scope=gpu-event-full-operator input=" << options.input
              << " reference=fp64-attention-to-fp32 atol=" << kAtol << " rtol=" << kRtol
              << " intermediate_allocated_bytes=" << intermediate_bytes
              << " intermediate_used_bytes=" << (version == "materialized" ? intermediate_bytes : 0)
              << " output_finite=" << (validation.finite ? "OK" : "FAIL")
              << " timed=" << (timing != nullptr) << " correct=" << (validation.correct ? "OK" : "FAIL")
              << " precheck=" << precheck << " postcheck=" << postcheck;
    if (timing)
        std::cout << " min_ms=" << timing->minimum_ms << " median_ms=" << timing->median_ms
                  << " mean_ms=" << timing->mean_ms;
    else
        std::cout << " min_ms=NA median_ms=NA mean_ms=NA";
    std::cout << " max_abs_error=" << validation.max_absolute_error
              << " max_tolerance_ratio=" << validation.max_tolerance_ratio << '\n';
}

void save_samples(const Options& options, const std::string& version, const Timing& timing) {
    if (options.samples.empty()) return;
    std::ofstream file(options.samples);
    if (!file) usage_error("cannot open sample file: " + options.samples);
    file << "config_id,implementation,runtime,shape,dtype,block,block_k,block_d,num_warps,launches,warmup,repeat,seed,process,scope,input,input_precision,atol,rtol,intermediate_allocated_bytes,intermediate_used_bytes,sample,ms\n";
    file << std::setprecision(17);
    const std::size_t intermediate_bytes = 2ULL * options.seq * options.seq * sizeof(float);
    for (std::size_t i = 0; i < timing.samples.size(); ++i)
        file << options.config_id << ",hip-" << version << ",hip,"
             << options.seq << 'x' << options.dim << ",float32," << options.block << ','
             << (version == "online" ? 1 : 0) << ",0,0," << (version == "materialized" ? 3 : 1)
             << ',' << options.warmup << ',' << options.repeat << ',' << options.seed << ','
             << options.process << ",gpu-event-full-operator," << options.input << ",fp32,"
             << kAtol << ',' << kRtol << ',' << intermediate_bytes << ','
             << (version == "materialized" ? intermediate_bytes : 0) << ',' << i << ',' << timing.samples[i] << '\n';
    if (!file) usage_error("failed writing sample file: " + options.samples);
}

int main(int argc, char** argv) {
    const Options options = parse_options(argc, argv);
    const int qkv_elements = options.seq * options.dim;
    const int score_elements = options.seq * options.seq;
    const std::size_t qkv_bytes = static_cast<std::size_t>(qkv_elements) * sizeof(float);
    const std::size_t score_bytes = static_cast<std::size_t>(score_elements) * sizeof(float);
    int device = 0;
    hipDeviceProp_t properties{};
    HIP_CHECK(hipGetDevice(&device));
    HIP_CHECK(hipGetDeviceProperties(&properties, device));
    std::cout << "ENV runtime=hip gpu=\"" << properties.name << "\" arch="
              << properties.gcnArchName << " wave_size=" << properties.warpSize << '\n';

    const Inputs input = make_input(options);
    const std::vector<float> reference = reference_attention(input, options.seq, options.dim);
    std::vector<float> actual(qkv_elements);
    float *d_q, *d_k, *d_value, *d_output, *d_scores, *d_probabilities;
    HIP_CHECK(hipMalloc(&d_q, qkv_bytes));
    HIP_CHECK(hipMalloc(&d_k, qkv_bytes));
    HIP_CHECK(hipMalloc(&d_value, qkv_bytes));
    HIP_CHECK(hipMalloc(&d_output, qkv_bytes));
    // Both paths keep the same host allocation plan. Online avoids accessing
    // scores/P, but this executable does not demonstrate a lower allocation peak.
    HIP_CHECK(hipMalloc(&d_scores, score_bytes));
    HIP_CHECK(hipMalloc(&d_probabilities, score_bytes));
    HIP_CHECK(hipMemcpy(d_q, input.q.data(), qkv_bytes, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(d_k, input.k.data(), qkv_bytes, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(d_value, input.value.data(), qkv_bytes, hipMemcpyHostToDevice));
    const float scale = static_cast<float>(1.0 / std::sqrt(static_cast<double>(options.dim)));
    const int score_grid = (score_elements + options.block - 1) / options.block;
    const int output_grid = (qkv_elements + options.block - 1) / options.block;
    const std::size_t online_shared = (options.block + options.dim + 4) * sizeof(float);
    auto launch = [&](const std::string& version) {
        if (version == "materialized") {
            hipLaunchKernelGGL(scores_kernel, dim3(score_grid), dim3(options.block), 0, 0,
                               d_q, d_k, d_scores, options.seq, options.dim, scale);
            HIP_CHECK(hipGetLastError());
            hipLaunchKernelGGL(softmax_rows_kernel, dim3(options.seq), dim3(options.block),
                               options.block * sizeof(float), 0,
                               d_scores, d_probabilities, options.seq);
            HIP_CHECK(hipGetLastError());
            hipLaunchKernelGGL(probability_value_kernel, dim3(output_grid), dim3(options.block), 0, 0,
                               d_probabilities, d_value, d_output, options.seq, options.dim);
        } else {
            hipLaunchKernelGGL(online_attention_kernel, dim3(options.seq), dim3(options.block),
                               online_shared, 0, d_q, d_k, d_value, d_output,
                               options.seq, options.dim, scale);
        }
        HIP_CHECK(hipGetLastError());
    };
    bool all_correct = true;
    for (const std::string version : {"materialized", "online"}) {
        if (options.version != "all" && options.version != version) continue;
        HIP_CHECK(hipMemset(d_output, 0xFF, qkv_bytes));
        HIP_CHECK(hipMemset(d_scores, 0xFF, score_bytes));
        HIP_CHECK(hipMemset(d_probabilities, 0xFF, score_bytes));
        launch(version);
        HIP_CHECK(hipDeviceSynchronize());
        HIP_CHECK(hipMemcpy(actual.data(), d_output, qkv_bytes, hipMemcpyDeviceToHost));
        const Validation precheck = validate(actual, reference);
        if (!precheck.correct) {
            print_result(options, version, precheck, nullptr, "FAIL", "NA");
            all_correct = false;
            continue;
        }
        const Timing timing = benchmark(options, [&] { launch(version); });
        HIP_CHECK(hipMemcpy(actual.data(), d_output, qkv_bytes, hipMemcpyDeviceToHost));
        const Validation postcheck = validate(actual, reference);
        print_result(options, version, postcheck, &timing, "OK", postcheck.correct ? "OK" : "FAIL");
        all_correct = all_correct && postcheck.correct;
        if (postcheck.correct) save_samples(options, version, timing);
    }
    HIP_CHECK(hipFree(d_q));
    HIP_CHECK(hipFree(d_k));
    HIP_CHECK(hipFree(d_value));
    HIP_CHECK(hipFree(d_output));
    HIP_CHECK(hipFree(d_scores));
    HIP_CHECK(hipFree(d_probabilities));
    return all_correct ? 0 : 1;
}

这份文件包含本路线的输入准备、kernel、启动、校验与计时；后续配置共用它，按函数和运行参数定位改动。

先编译一次。未修改源码时可复用本轮可执行文件；修改源码后请新建结果目录再编译：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
if [[ ! -x "$RUN/build/attention_hip" ]]; then
  hipcc --offload-arch=gfx1201 -O3 -std=c++17 \
    chapter12/attention_hip.hip -o "$RUN/build/attention_hip"
fi

launch() 用三个 kernel 对应三个阶段：

阶段	分工	写出什么
scores_kernel	一个线程计算一个 query-key 点积	一个 Score
softmax_rows_kernel	一个 block 协作处理一行分数	这一行概率
probability_value_kernel	一个线程累加一个输出分量	一个 O 元素

同一 query 的不同 key 点积互不依赖，分数阶段可以同时展开多个 query-key 对；最终加权和则沿 key 归约。行 Softmax 的 LDS 复用遵循第 10 章的同步规则：归约完成后，必须等所有线程读完最大值，才可用同一片 LDS 保存下一次归约的部分和。

先运行物化路径，为本路线建立基线：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
test ! -e "$RUN/n128x64-hip-materialized-b256-p1.csv" && \
"$RUN/build/attention_hip" \
  --block 256 --version materialized --seq 128 \
  --dim 64 --warmup 10 --repeat 50 \
  --seed 20260920 --input dyadic --config-id hip-materialized-b256 \
  --samples "$RUN/n128x64-hip-materialized-b256-p1.csv" --process 1

真实输出：hip-materialized-b256，shape=128x64，第 1 个独立进程

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 wave_size=32
RESULT operator=attention implementation=hip-materialized runtime=hip config_id=hip-materialized-b256 shape=128x64 seq=128 dim=64 dtype=float32 input_precision=fp32 block=256 block_k=0 block_d=0 num_warps=0 launches=3 warmup=10 repeat=50 seed=20260920 process=1 timing=gpu-event scope=gpu-event-full-operator input=dyadic reference=fp64-attention-to-fp32 atol=2.0000000000000002e-05 rtol=0.00020000000000000001 intermediate_allocated_bytes=131072 intermediate_used_bytes=131072 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.059241000562906265 median_ms=0.059581000357866287 mean_ms=0.059687860012054443 max_abs_error=2.2351741790771484e-08 max_tolerance_ratio=0.00098910493985020415

先检查 correct/precheck/postcheck、有限值与容差比，再核对 shape=128x64,launches=3。主形状基线完整时间为 59.681 μs，三个进程中位数范围为 59.581–59.741 μs。源码中的 Scores/P 交接已经提供了融合假设，不需要先采 profiler 才能提出它。

这条路径保存了完整中间矩阵，也进行了三次 kernel 调用。接下来尝试在一个 block 内直接累计一行输出，检验这些中间成本是否值得省去。

12.5.2 HIP 改动：一个 block 完成在线累计

在线版本用 blockIdx.x 选择 query 行。每读一个 key，线程协作算完点积，再更新这一行的最大值、指数和与向量累计。这样最终只需写出 O[S,D]，kernel 不再读写 Scores/P。

线程 0 更新当前 key 的标量状态：

if (tid == 0) {
    float score = reduction[0] * scale;
    float old_max = state[0];
    float new_max = fmaxf(old_max, score);
    float alpha = expf(old_max - new_max);
    float beta = expf(score - new_max);
    state[0] = new_max;
    state[1] = state[1] * alpha + beta;
    state[2] = alpha;
    state[3] = beta;
}

屏障之后，其他线程读取本轮的 alpha/beta，分别更新自己负责的 numerator[d]。更新结束后再同步，下一轮才能覆盖系数。

HIP · 固定 key 2 的共享依赖系数写好、取得副本、消费完成

1/5

固定已经算完 score=4 的时刻旧最大值为 2、l≈1.3679、numerator≈[1,0.7358]，都保存在共享空间。本图不重演前面的点积归约。

线程 0 更新共享状态与系数线程 0 写 m=4、l≈1.1851、α=exp(−2)、β=1。写好不等于其他线程已取得这些值，numerator 尚未改变。

全 block 到达屏障，消费者取得系数负责两个分量的线程现在可读取同一组 α/β，保留自己的副本。共享 numerator 仍是旧值。

分别更新共享 numerator 的两项每个消费者读取自己的旧 a[d]，计算 α×旧 a[d]+β×V₂[d]，分别写成 3.1353 和 1.0996。

读者完成使用后才允许覆盖全 block 再次同步，保护仍在消费当前系数的线程。只有消费完成后，后续迭代才可复用共享位置；本例 key 2 已是最后一项。

固定已经算完 score=4 的时刻旧最大值为 2、l≈1.3679、numerator≈[1,0.7358]，都保存在共享空间。本图不重演前面的点积归约。

步骤 1 / 5旧状态

只放大两个输出分量的消费者；其余线程仍参与协作。这里不统计全部屏障，也不表示执行耗时。

图 12.4　只放大 key 2 的系数交接，起点是已经得到 score=4。线程 0 写好共享系数，全 block 同步后消费者才能取得它们；各自更新共享 numerator 后，还需等待消费完成才能复用系数。图中仅展开两个输出分量，其余线程仍参与协作，不统计完整 kernel 的屏障次数。

图 12.4 标出了系数写入与消费之间的依赖。物化基线会展开多个 query-key 对，在线版则在每个 block 内按顺序遍历 key；它还需要保存跨轮状态，并多次等待点积和状态更新完成。省去中间矩阵，并不意味着剩余计算的分工与协作成本保持不变。

在线算法改变了跨轮状态，先让后续 key 的最大值持续升高，检查分母与分子是否一起重缩放：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
"$RUN/build/attention_hip" \
  --block 256 --version online --seq 65 \
  --dim 64 --warmup 0 --repeat 1 \
  --seed 20260920 --input rising-max --config-id hip-online-b256

正确性输出：HIP online，65×64，rising-max

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 wave_size=32
RESULT operator=attention implementation=hip-online runtime=hip config_id=hip-online-b256 shape=65x64 seq=65 dim=64 dtype=float32 input_precision=fp32 block=256 block_k=1 block_d=0 num_warps=0 launches=1 warmup=0 repeat=1 seed=20260920 process=1 timing=gpu-event scope=gpu-event-full-operator input=rising-max reference=fp64-attention-to-fp32 atol=2.0000000000000002e-05 rtol=0.00020000000000000001 intermediate_allocated_bytes=33800 intermediate_used_bytes=0 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.12640099227428436 median_ms=0.12640099227428436 mean_ms=0.12640099227428436 max_abs_error=1.4901161193847656e-08 max_tolerance_ratio=0.00060550631823742102

再检查只有一个 key 的情形，输出应等于 value：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
"$RUN/build/attention_hip" \
  --block 256 --version online --seq 1 \
  --dim 1 --warmup 0 --repeat 1 \
  --seed 20260920 --input dyadic --config-id hip-online-b256

正确性输出：HIP online，单 key

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 wave_size=32
RESULT operator=attention implementation=hip-online runtime=hip config_id=hip-online-b256 shape=1x1 seq=1 dim=1 dtype=float32 input_precision=fp32 block=256 block_k=1 block_d=0 num_warps=0 launches=1 warmup=0 repeat=1 seed=20260920 process=1 timing=gpu-event scope=gpu-event-full-operator input=dyadic reference=fp64-attention-to-fp32 atol=2.0000000000000002e-05 rtol=0.00020000000000000001 intermediate_allocated_bytes=8 intermediate_used_bytes=0 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.012880000285804272 median_ms=0.012880000285804272 mean_ms=0.012880000285804272 max_abs_error=0 max_tolerance_ratio=0

两条命令仅读正确性，不比较一次采样的时间。检查通过后，回到与基线相同的 dyadic 主输入，测完整在线版本：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
test ! -e "$RUN/n128x64-hip-online-b256-p1.csv" && \
"$RUN/build/attention_hip" \
  --block 256 --version online --seq 128 \
  --dim 64 --warmup 10 --repeat 50 \
  --seed 20260920 --input dyadic --config-id hip-online-b256 \
  --samples "$RUN/n128x64-hip-online-b256-p1.csv" --process 1

真实输出：hip-online-b256，shape=128x64，第 1 个独立进程

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 wave_size=32
RESULT operator=attention implementation=hip-online runtime=hip config_id=hip-online-b256 shape=128x64 seq=128 dim=64 dtype=float32 input_precision=fp32 block=256 block_k=1 block_d=0 num_warps=0 launches=1 warmup=10 repeat=50 seed=20260920 process=1 timing=gpu-event scope=gpu-event-full-operator input=dyadic reference=fp64-attention-to-fp32 atol=2.0000000000000002e-05 rtol=0.00020000000000000001 intermediate_allocated_bytes=131072 intermediate_used_bytes=0 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.22992299497127533 median_ms=0.23068299889564514 mean_ms=0.23063763827085496 max_abs_error=2.1653249859809875e-08 max_tolerance_ratio=0.0010513768060094045

先核对检查通过、launches=1 与 intermediate_used_bytes=0，再读完整时间；后者不代表主机没有分配这两块缓冲。

图 12.5　S=128、D=64，两版 block 都为 256，比较完整前向路径。在线版减少目标 kernel 数和中间读写，但没有保持相同的 key 并行分工。

在线版为 230.643 μs，范围 230.523–230.683 μs，用时约为物化基线的 3.86 倍。独立确认仍为 230.602 μs 对 59.961 μs。本轮保留物化实现作为 HIP 方案；在线公式通过了正确性检查，但这套执行安排没有带来性能收益。

这个负结果修正了最初的假设：仅仅省去中间矩阵还不够，还要考虑原本可以展开的 key 工作，现在怎样被顺序处理。

12.5.3 HIP 诊断：区分阶段时间与原因

这次要回答：完整路径的退步，是否也出现在在线 kernel 的设备执行区间？ 这会帮助决定继续研究设备端的执行安排，还是先调查测量与提交方式。仍在 code/part2-kernels/，分别采两条路径；每个 kernel 按名字筛选：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
mkdir "$RUN/profile-hip-materialized-b256" && \
rocprofv3 --kernel-trace --output-format csv \
  --output-directory "$RUN/profile-hip-materialized-b256" --output-file hip-materialized-b256 -- \
  "$RUN/build/attention_hip" \
    --block 256 --version materialized --seq 128 \
    --dim 64 --warmup 5 --repeat 10 \
    --seed 20260920 --input dyadic --config-id hip-materialized-b256 && \
python chapter8/inspect_trace.py \
  "$RUN/profile-hip-materialized-b256/hip-materialized-b256_kernel_trace.csv" \
  --kernel scores_kernel --skip 6 --take 10 && \
python chapter8/inspect_trace.py \
  "$RUN/profile-hip-materialized-b256/hip-materialized-b256_kernel_trace.csv" \
  --kernel softmax_rows_kernel --skip 6 --take 10 && \
python chapter8/inspect_trace.py \
  "$RUN/profile-hip-materialized-b256/hip-materialized-b256_kernel_trace.csv" \
  --kernel probability_value_kernel --skip 6 --take 10

筛选后的真实 trace：hip-materialized-b256，shape=128x64

kernel_name="scores_kernel(float const*, float const*, float*, int, int, float)"
csv_rows=51 target_rows=16 skipped_rows=6 kept_rows=10
sort=Start_Timestamp_ascending
first_start_ns=887497031786019 first_end_ns=887497031805419 first_duration_ns=19400
kernel_min_us=8.040000 kernel_median_us=8.140500 kernel_max_us=19.400000
grid_work_items_xyz=16384x1x1 workgroup_size_xyz=256x1x1 workgroups_xyz=64x1x1
workgroups_total=64
vgpr_count=8 sgpr_count=128 lds_bytes=0 scratch_bytes=0

kernel_name="softmax_rows_kernel(float const*, float*, int)"
csv_rows=51 target_rows=16 skipped_rows=6 kept_rows=10
sort=Start_Timestamp_ascending
first_start_ns=887497031809699 first_end_ns=887497031814260 first_duration_ns=4561
kernel_min_us=4.480000 kernel_median_us=4.560000 kernel_max_us=4.640000
grid_work_items_xyz=32768x1x1 workgroup_size_xyz=256x1x1 workgroups_xyz=128x1x1
workgroups_total=128
vgpr_count=16 sgpr_count=128 lds_bytes=0 scratch_bytes=0

kernel_name="probability_value_kernel(float const*, float const*, float*, int, int)"
csv_rows=51 target_rows=16 skipped_rows=6 kept_rows=10
sort=Start_Timestamp_ascending
first_start_ns=887497031845800 first_end_ns=887497031880681 first_duration_ns=34881
kernel_min_us=34.761000 kernel_median_us=34.921000 kernel_max_us=35.040000
grid_work_items_xyz=8192x1x1 workgroup_size_xyz=256x1x1 workgroups_xyz=32x1x1
workgroups_total=32
vgpr_count=16 sgpr_count=128 lds_bytes=0 scratch_bytes=0

原始 kernel trace CSV：hip-materialized-b256，shape=128x64

"Kind","Agent_Id","Queue_Id","Stream_Id","Thread_Id","Dispatch_Id","Kernel_Id","Kernel_Name","Correlation_Id","Start_Timestamp","End_Timestamp","LDS_Block_Size","Scratch_Size","VGPR_Count","Accum_VGPR_Count","SGPR_Count","Workgroup_Size_X","Workgroup_Size_Y","Workgroup_Size_Z","Grid_Size_X","Grid_Size_Y","Grid_Size_Z"
"KERNEL_DISPATCH","Agent 1",1,0,1226358,1,13,"__amd_rocclr_fillBufferUnAligned",1,887497019338020,887497019342420,0,0,8,0,128,256,1,1,2048,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,2,13,"__amd_rocclr_fillBufferUnAligned",2,887497019527222,887497019529382,0,0,8,0,128,256,1,1,4096,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,3,13,"__amd_rocclr_fillBufferUnAligned",3,887497019534822,887497019537502,0,0,8,0,128,256,1,1,4096,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,4,22,"scores_kernel(float const*, float const*, float*, int, int, float)",4,887497031213151,887497031221671,0,0,8,0,128,256,1,1,16384,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,6,20,"probability_value_kernel(float const*, float const*, float*, int, int)",6,887497031235871,887497031274031,0,0,16,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,5,21,"softmax_rows_kernel(float const*, float*, int)",5,887497031226351,887497031231351,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,7,22,"scores_kernel(float const*, float const*, float*, int, int, float)",7,887497031366783,887497031385263,0,0,8,0,128,256,1,1,16384,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,8,21,"softmax_rows_kernel(float const*, float*, int)",8,887497031389743,887497031394383,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,9,20,"probability_value_kernel(float const*, float const*, float*, int, int)",9,887497031412503,887497031450264,0,0,16,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,10,22,"scores_kernel(float const*, float const*, float*, int, int, float)",10,887497031454704,887497031462864,0,0,8,0,128,256,1,1,16384,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,11,21,"softmax_rows_kernel(float const*, float*, int)",11,887497031467224,887497031484384,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,12,20,"probability_value_kernel(float const*, float const*, float*, int, int)",12,887497031488824,887497031526665,0,0,16,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,13,22,"scores_kernel(float const*, float const*, float*, int, int, float)",13,887497031531105,887497031539345,0,0,8,0,128,256,1,1,16384,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,14,21,"softmax_rows_kernel(float const*, float*, int)",14,887497031543825,887497031558706,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,15,20,"probability_value_kernel(float const*, float const*, float*, int, int)",15,887497031563146,887497031599186,0,0,16,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,16,22,"scores_kernel(float const*, float const*, float*, int, int, float)",16,887497031603386,887497031611546,0,0,8,0,128,256,1,1,16384,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,17,21,"softmax_rows_kernel(float const*, float*, int)",17,887497031615786,887497031632627,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,18,20,"probability_value_kernel(float const*, float const*, float*, int, int)",18,887497031636947,887497031671907,0,0,16,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,19,22,"scores_kernel(float const*, float const*, float*, int, int, float)",19,887497031676147,887497031684267,0,0,8,0,128,256,1,1,16384,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,20,21,"softmax_rows_kernel(float const*, float*, int)",20,887497031688428,887497031693148,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,21,20,"probability_value_kernel(float const*, float const*, float*, int, int)",21,887497031710188,887497031745268,0,0,16,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,22,22,"scores_kernel(float const*, float const*, float*, int, int, float)",22,887497031786019,887497031805419,0,0,8,0,128,256,1,1,16384,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,23,21,"softmax_rows_kernel(float const*, float*, int)",23,887497031809699,887497031814260,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,24,20,"probability_value_kernel(float const*, float const*, float*, int, int)",24,887497031845800,887497031880681,0,0,16,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,25,22,"scores_kernel(float const*, float const*, float*, int, int, float)",25,887497031916561,887497031934041,0,0,8,0,128,256,1,1,16384,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,26,21,"softmax_rows_kernel(float const*, float*, int)",26,887497031938281,887497031942841,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,27,20,"probability_value_kernel(float const*, float const*, float*, int, int)",27,887497031961122,887497031996082,0,0,16,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,28,22,"scores_kernel(float const*, float const*, float*, int, int, float)",28,887497032026443,887497032034683,0,0,8,0,128,256,1,1,16384,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,29,21,"softmax_rows_kernel(float const*, float*, int)",29,887497032051403,887497032055883,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,30,20,"probability_value_kernel(float const*, float const*, float*, int, int)",30,887497032078044,887497032113004,0,0,16,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,31,22,"scores_kernel(float const*, float const*, float*, int, int, float)",31,887497032142405,887497032150525,0,0,8,0,128,256,1,1,16384,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,32,21,"softmax_rows_kernel(float const*, float*, int)",32,887497032166525,887497032171085,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,33,20,"probability_value_kernel(float const*, float const*, float*, int, int)",33,887497032197535,887497032232296,0,0,16,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,34,22,"scores_kernel(float const*, float const*, float*, int, int, float)",34,887497032261136,887497032269257,0,0,8,0,128,256,1,1,16384,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,35,21,"softmax_rows_kernel(float const*, float*, int)",35,887497032285577,887497032290217,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,36,20,"probability_value_kernel(float const*, float const*, float*, int, int)",36,887497032316177,887497032351098,0,0,16,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,37,22,"scores_kernel(float const*, float const*, float*, int, int, float)",37,887497032380338,887497032388418,0,0,8,0,128,256,1,1,16384,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,38,21,"softmax_rows_kernel(float const*, float*, int)",38,887497032404539,887497032409139,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,39,20,"probability_value_kernel(float const*, float const*, float*, int, int)",39,887497032436719,887497032471640,0,0,16,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,40,22,"scores_kernel(float const*, float const*, float*, int, int, float)",40,887497032501280,887497032509480,0,0,8,0,128,256,1,1,16384,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,41,21,"softmax_rows_kernel(float const*, float*, int)",41,887497032526001,887497032530521,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,42,20,"probability_value_kernel(float const*, float const*, float*, int, int)",42,887497032556701,887497032591502,0,0,16,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,43,22,"scores_kernel(float const*, float const*, float*, int, int, float)",43,887497032620342,887497032628422,0,0,8,0,128,256,1,1,16384,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,44,21,"softmax_rows_kernel(float const*, float*, int)",44,887497032644942,887497032649502,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,45,20,"probability_value_kernel(float const*, float const*, float*, int, int)",45,887497032675313,887497032710353,0,0,16,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,46,22,"scores_kernel(float const*, float const*, float*, int, int, float)",46,887497032740154,887497032748314,0,0,8,0,128,256,1,1,16384,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,47,21,"softmax_rows_kernel(float const*, float*, int)",47,887497032764874,887497032769394,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,48,20,"probability_value_kernel(float const*, float const*, float*, int, int)",48,887497032793595,887497032828555,0,0,16,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,49,22,"scores_kernel(float const*, float const*, float*, int, int, float)",49,887497032857916,887497032865956,0,0,8,0,128,256,1,1,16384,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,50,21,"softmax_rows_kernel(float const*, float*, int)",50,887497032882356,887497032886876,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226358,51,20,"probability_value_kernel(float const*, float const*, float*, int, int)",51,887497032912706,887497032947627,0,0,16,0,128,256,1,1,8192,1,1

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
mkdir "$RUN/profile-hip-online-b256" && \
rocprofv3 --kernel-trace --output-format csv \
  --output-directory "$RUN/profile-hip-online-b256" --output-file hip-online-b256 -- \
  "$RUN/build/attention_hip" \
    --block 256 --version online --seq 128 \
    --dim 64 --warmup 5 --repeat 10 \
    --seed 20260920 --input dyadic --config-id hip-online-b256 && \
python chapter8/inspect_trace.py \
  "$RUN/profile-hip-online-b256/hip-online-b256_kernel_trace.csv" \
  --kernel online_attention_kernel --skip 6 --take 10

筛选后的真实 trace：hip-online-b256，shape=128x64

kernel_name="online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)"
csv_rows=19 target_rows=16 skipped_rows=6 kept_rows=10
sort=Start_Timestamp_ascending
first_start_ns=887497198778368 first_end_ns=887497199001372 first_duration_ns=223004
kernel_min_us=222.683000 kernel_median_us=222.903000 kernel_max_us=240.603000
grid_work_items_xyz=32768x1x1 workgroup_size_xyz=256x1x1 workgroups_xyz=128x1x1
workgroups_total=128
vgpr_count=16 sgpr_count=128 lds_bytes=0 scratch_bytes=0

原始 kernel trace CSV：hip-online-b256，shape=128x64

"Kind","Agent_Id","Queue_Id","Stream_Id","Thread_Id","Dispatch_Id","Kernel_Id","Kernel_Name","Correlation_Id","Start_Timestamp","End_Timestamp","LDS_Block_Size","Scratch_Size","VGPR_Count","Accum_VGPR_Count","SGPR_Count","Workgroup_Size_X","Workgroup_Size_Y","Workgroup_Size_Z","Grid_Size_X","Grid_Size_Y","Grid_Size_Z"
"KERNEL_DISPATCH","Agent 1",1,0,1226368,1,7,"__amd_rocclr_fillBufferUnAligned",1,887497185858205,887497185863325,0,0,8,0,128,256,1,1,2048,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,3,7,"__amd_rocclr_fillBufferUnAligned",3,887497186061257,887497186063977,0,0,8,0,128,256,1,1,4096,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,2,7,"__amd_rocclr_fillBufferUnAligned",2,887497186053457,887497186055617,0,0,8,0,128,256,1,1,4096,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,4,19,"online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)",4,887497197260764,887497197502048,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,5,19,"online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)",5,887497197602929,887497197827973,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,6,19,"online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)",6,887497197832453,887497198055896,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,7,19,"online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)",7,887497198060216,887497198283460,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,8,19,"online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)",8,887497198287620,887497198510664,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,9,19,"online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)",9,887497198515024,887497198738147,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,10,19,"online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)",10,887497198778368,887497199001372,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,11,19,"online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)",11,887497199039912,887497199262956,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,12,19,"online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)",12,887497199292596,887497199515720,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,13,19,"online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)",13,887497199545600,887497199768324,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,14,19,"online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)",14,887497199798325,887497200021128,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,15,19,"online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)",15,887497200051049,887497200273892,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,16,19,"online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)",16,887497200303453,887497200526416,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,17,19,"online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)",17,887497200555537,887497200778340,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,18,19,"online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)",18,887497200808061,887497201030744,0,0,16,0,128,256,1,1,32768,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1226368,19,19,"online_attention_kernel(float const*, float const*, float const*, float*, int, int, float)",19,887497212227815,887497212468418,0,0,16,0,128,256,1,1,32768,1,1

先读报告的 target_rows/skipped_rows/kept_rows，核对分析窗口；再读名称与 launch，最后比较 kernel_median_us 和范围。物化三个阶段分别约 8.141、4.560、34.921 μs，在线则如下文。不要把三个独立中位数相加当成完整 benchmark。

先在 trace 中按 scores_kernel、softmax_rows_kernel、probability_value_kernel 筛选物化路径，再查看 online_attention_kernel。确认同一 shape、block 参数和完整调用范围，排除初始化与预检调用。阶段 trace 能显示耗时主要落在哪个目标 kernel，独立采集时的总记录数不是一次 Attention 的调用数。

本次每个目标 kernel 都有 16 条记录：1 次预检、5 次预热、10 次目标重复，排除前 6 条后统计。在线 kernel 的 trace 中位数为 222.903 μs，说明较长的设备执行本身值得继续研究；不能拿它与另一轮 event 相减，计算某项启动或同步成本。

本轮 block=256,D=64 时，Softmax 动态 LDS 为 256×4=1024 B，在线版为 (256+64+4)×4=1296 B。这些容量来自实际启动参数。trace 的 LDS_Block_Size 为 0，没有包含这项动态申请，因此不能读作“不使用 LDS”。

源码确认在线版在每个 key 后需要协作，trace 确认它的执行时间；两者还不能说明每个屏障贡献了多少延迟。继续优化时，可以提出“按 key tile 合并多轮工作”的下一项假设，但必须保留相同公式与正确性检查，再建立单独对照。直接删掉承担数据依赖的屏障，会先改变程序是否正确。

12.5.4 Triton 基线：一行 query 的三个阶段

完整源码：code/part2-kernels/chapter12/attention_triton.py

"""FP32 materialized and online Attention for Chapter 12.

Single head, non-causal, Q/K/V shape [SEQ, DIM]. Every GPU multiplication and
reduction stays FP32; no tl.dot or reduced-precision tensor-core path is used.
Timing covers all kernels of one forward call, excluding allocation/reference.
"""

from __future__ import annotations

import argparse
import csv
import math
import statistics
from dataclasses import dataclass
from pathlib import Path

import torch
import triton
import triton.language as tl


ATOL = 2.0e-5
RTOL = 2.0e-4
MATERIALIZED_BLOCK_K = 32
NUM_WARPS = 4


@triton.jit
def materialized_scores_kernel(
    q_ptr, k_ptr, scores_ptr, scale,
    SEQ: tl.constexpr, DIM: tl.constexpr,
    BLOCK_K: tl.constexpr, BLOCK_D: tl.constexpr,
):
    row = tl.program_id(0)
    d = tl.arange(0, BLOCK_D)
    q = tl.load(q_ptr + row * DIM + d, mask=d < DIM, other=0.0)
    for key_start in range(0, SEQ, BLOCK_K):
        keys = key_start + tl.arange(0, BLOCK_K)
        key_mask = keys < SEQ
        k = tl.load(
            k_ptr + keys[:, None] * DIM + d[None, :],
            mask=key_mask[:, None] & (d[None, :] < DIM), other=0.0,
        )
        scores = tl.sum(k * q[None, :], axis=1) * scale
        tl.store(scores_ptr + row * SEQ + keys, scores, mask=key_mask)


@triton.jit
def materialized_softmax_kernel(
    scores_ptr, probabilities_ptr,
    SEQ: tl.constexpr, BLOCK_S: tl.constexpr,
):
    row = tl.program_id(0)
    keys = tl.arange(0, BLOCK_S)
    scores = tl.load(scores_ptr + row * SEQ + keys, mask=keys < SEQ, other=-float("inf"))
    weights = tl.exp(scores - tl.max(scores, axis=0))
    probabilities = weights / tl.sum(weights, axis=0)
    tl.store(probabilities_ptr + row * SEQ + keys, probabilities, mask=keys < SEQ)


@triton.jit
def materialized_pv_kernel(
    probabilities_ptr, v_ptr, output_ptr,
    SEQ: tl.constexpr, DIM: tl.constexpr,
    BLOCK_K: tl.constexpr, BLOCK_D: tl.constexpr,
):
    row = tl.program_id(0)
    d = tl.arange(0, BLOCK_D)
    accumulator = tl.zeros((BLOCK_D,), tl.float32)
    for key_start in range(0, SEQ, BLOCK_K):
        keys = key_start + tl.arange(0, BLOCK_K)
        key_mask = keys < SEQ
        probabilities = tl.load(probabilities_ptr + row * SEQ + keys, mask=key_mask, other=0.0)
        values = tl.load(
            v_ptr + keys[:, None] * DIM + d[None, :],
            mask=key_mask[:, None] & (d[None, :] < DIM), other=0.0,
        )
        accumulator += tl.sum(probabilities[:, None] * values, axis=0)
    tl.store(output_ptr + row * DIM + d, accumulator, mask=d < DIM)


@triton.jit
def online_attention_kernel(
    q_ptr, k_ptr, v_ptr, output_ptr, scale,
    SEQ: tl.constexpr, DIM: tl.constexpr,
    BLOCK_K: tl.constexpr, BLOCK_D: tl.constexpr,
):
    row = tl.program_id(0)
    d = tl.arange(0, BLOCK_D)
    q = tl.load(q_ptr + row * DIM + d, mask=d < DIM, other=0.0)
    running_max = -float("inf")
    running_sum = 0.0
    accumulator = tl.zeros((BLOCK_D,), tl.float32)
    for key_start in range(0, SEQ, BLOCK_K):
        keys = key_start + tl.arange(0, BLOCK_K)
        key_mask = keys < SEQ
        k = tl.load(
            k_ptr + keys[:, None] * DIM + d[None, :],
            mask=key_mask[:, None] & (d[None, :] < DIM), other=0.0,
        )
        scores = tl.sum(k * q[None, :], axis=1) * scale
        scores = tl.where(key_mask, scores, -float("inf"))
        # SEQ > 0 and key_start < SEQ guarantee a valid key in every tile.
        tile_max = tl.max(scores, axis=0)
        next_max = tl.maximum(running_max, tile_max)
        history_scale = tl.exp(running_max - next_max)
        probabilities = tl.exp(scores - next_max)
        values = tl.load(
            v_ptr + keys[:, None] * DIM + d[None, :],
            mask=key_mask[:, None] & (d[None, :] < DIM), other=0.0,
        )
        accumulator = accumulator * history_scale + tl.sum(
            probabilities[:, None] * values, axis=0
        )
        running_sum = running_sum * history_scale + tl.sum(probabilities, axis=0)
        running_max = next_max
    tl.store(output_ptr + row * DIM + d, accumulator / running_sum, mask=d < DIM)


@dataclass(frozen=True)
class Implementation:
    version: str
    block_k: int

    @property
    def name(self) -> str:
        return f"triton-{self.version}"

    @property
    def launches(self) -> int:
        return 3 if self.version == "materialized" else 1


@dataclass(frozen=True)
class Validation:
    correct: bool
    finite: bool
    max_absolute_error: float
    max_tolerance_ratio: float


@dataclass(frozen=True)
class Timing:
    minimum_ms: float
    median_ms: float
    mean_ms: float
    samples: tuple[float, ...]


def parse_args() -> argparse.Namespace:
    parser = argparse.ArgumentParser(description=__doc__)
    parser.add_argument("--version", choices=("all", "materialized", "online", "t0", "t1"), default="all")
    parser.add_argument("--seq", type=int, default=128)
    parser.add_argument("--dim", type=int, default=64)
    parser.add_argument("--block-k", type=int, choices=(16, 32, 64), default=None,
                        help="online key tile; default 32 (t0/t1 aliases select 16/32)")
    parser.add_argument("--num-warps", type=int, choices=(4,), default=4)
    parser.add_argument("--input", choices=("dyadic", "rising-max"), default="dyadic")
    parser.add_argument("--warmup", type=int, default=10)
    parser.add_argument("--repeat", type=int, default=50)
    parser.add_argument("--seed", type=int, default=20260920)
    parser.add_argument("--samples", type=Path)
    parser.add_argument("--config-id", default="")
    parser.add_argument("--process", type=int, default=1)
    args = parser.parse_args()
    if not 1 <= args.seq <= 4096 or not 1 <= args.dim <= 256:
        parser.error("require 1 <= seq <= 4096 and 1 <= dim <= 256")
    if args.warmup < 0 or args.repeat <= 0 or args.process <= 0:
        parser.error("warmup must be non-negative; repeat/process must be positive")
    if not 0 <= args.seed <= 0xFFFFFFFF:
        parser.error("seed must fit uint32")
    if any(c not in "abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789-_." for c in args.config_id):
        parser.error("invalid config-id")
    if args.samples and (args.version == "all" or not args.config_id):
        parser.error("--samples requires one version and a --config-id")
    if args.version in ("t0", "t1"):
        alias_tile = 16 if args.version == "t0" else 32
        if args.block_k is not None and args.block_k != alias_tile:
            parser.error("--block-k conflicts with the t0/t1 alias")
        args.block_k = alias_tile
        args.version = "online"
    if args.version == "materialized" and args.block_k not in (None, MATERIALIZED_BLOCK_K):
        parser.error("materialized uses a fixed key tile of 32")
    args.block_k = args.block_k or 32
    return args


def selected_implementations(args: argparse.Namespace) -> list[Implementation]:
    materialized = Implementation("materialized", MATERIALIZED_BLOCK_K)
    online = Implementation("online", args.block_k)
    if args.version == "all":
        return [materialized, online]
    return [materialized if args.version == "materialized" else online]


def make_host_input(seq: int, dim: int, seed: int, kind: str) -> tuple[torch.Tensor, ...]:
    """Integer formulas match attention_hip.hip exactly before FP32 division."""
    rows = torch.arange(seq, dtype=torch.int64).view(seq, 1)
    columns = torch.arange(dim, dtype=torch.int64).view(1, dim)
    q = ((17 * rows + 13 * columns + seed) % 65 - 32).float() / 64.0
    k = ((29 * rows + 7 * columns + seed + 11) % 65 - 32).float() / 64.0
    v = ((11 * rows + 19 * columns + seed + 23) % 65 - 32).float() / 32.0
    if kind == "rising-max":
        q.zero_()
        k.zero_()
        q[:, 0] = 8.0
        k[:, 0] = (torch.arange(seq, dtype=torch.float32) - seq // 2) / 8.0
    return q, k, v


def cpu_reference(q: torch.Tensor, k: torch.Tensor, v: torch.Tensor) -> torch.Tensor:
    """All reference operations run on CPU in FP64, then round once to FP32."""
    scores = (q.double() @ k.double().T) / math.sqrt(q.shape[1])
    return (torch.softmax(scores, dim=1) @ v.double()).float()


def launch(implementation: Implementation, q: torch.Tensor, k: torch.Tensor,
           v: torch.Tensor, output: torch.Tensor,
           scores: torch.Tensor, probabilities: torch.Tensor) -> None:
    seq, dim = q.shape
    block_d = triton.next_power_of_2(dim)
    parameters = dict(SEQ=seq, DIM=dim, BLOCK_K=implementation.block_k,
                      BLOCK_D=block_d, num_warps=NUM_WARPS)
    if implementation.version == "materialized":
        materialized_scores_kernel[(seq,)](q, k, scores, 1.0 / math.sqrt(dim), **parameters)
        materialized_softmax_kernel[(seq,)](
            scores, probabilities, SEQ=seq, BLOCK_S=triton.next_power_of_2(seq), num_warps=NUM_WARPS,
        )
        materialized_pv_kernel[(seq,)](probabilities, v, output, **parameters)
    else:
        online_attention_kernel[(seq,)](q, k, v, output, 1.0 / math.sqrt(dim), **parameters)


def validate(output: torch.Tensor, reference: torch.Tensor) -> Validation:
    actual = output.detach().to(device="cpu")
    finite = bool(torch.isfinite(actual).all() and torch.isfinite(reference).all())
    if not finite:
        return Validation(False, False, math.inf, math.inf)
    error = (actual.double() - reference.double()).abs()
    tolerance = ATOL + RTOL * reference.double().abs()
    return Validation(bool((error <= tolerance).all()), True,
                      float(error.max()), float((error / tolerance).max()))


def time_launch(implementation: Implementation, args: argparse.Namespace,
                q: torch.Tensor, k: torch.Tensor, v: torch.Tensor,
                output: torch.Tensor, scores: torch.Tensor, probabilities: torch.Tensor) -> Timing:
    for _ in range(args.warmup):
        launch(implementation, q, k, v, output, scores, probabilities)
    torch.cuda.synchronize()
    start = torch.cuda.Event(enable_timing=True)
    stop = torch.cuda.Event(enable_timing=True)
    times: list[float] = []
    for _ in range(args.repeat):
        start.record()
        launch(implementation, q, k, v, output, scores, probabilities)
        stop.record()
        stop.synchronize()
        times.append(float(start.elapsed_time(stop)))
    return Timing(min(times), statistics.median(times), statistics.fmean(times), tuple(times))


def metadata(implementation: Implementation, args: argparse.Namespace) -> dict[str, object]:
    intermediate_bytes = 2 * args.seq * args.seq * 4
    return dict(
        config_id=args.config_id, implementation=implementation.name, runtime="triton",
        shape=f"{args.seq}x{args.dim}", dtype="float32", block=0,
        block_k=implementation.block_k, block_d=triton.next_power_of_2(args.dim),
        num_warps=NUM_WARPS, launches=implementation.launches,
        warmup=args.warmup, repeat=args.repeat, seed=args.seed, process=args.process,
        scope="gpu-event-full-operator", input=args.input, input_precision="fp32",
        atol=ATOL, rtol=RTOL, intermediate_allocated_bytes=intermediate_bytes,
        intermediate_used_bytes=intermediate_bytes if implementation.version == "materialized" else 0,
    )


def print_result(implementation: Implementation, args: argparse.Namespace,
                 validation: Validation, timing: Timing | None,
                 precheck: str, postcheck: str) -> None:
    fields = metadata(implementation, args)
    fields.update(seq=args.seq, dim=args.dim, programs_per_query=1, timing="gpu-event",
                  reference="fp64-attention-to-fp32", output_finite="OK" if validation.finite else "FAIL",
                  timed=int(timing is not None), correct="OK" if validation.correct else "FAIL",
                  precheck=precheck, postcheck=postcheck,
                  min_ms="NA" if timing is None else format(timing.minimum_ms, ".17g"),
                  median_ms="NA" if timing is None else format(timing.median_ms, ".17g"),
                  mean_ms="NA" if timing is None else format(timing.mean_ms, ".17g"),
                  max_abs_error=validation.max_absolute_error,
                  max_tolerance_ratio=validation.max_tolerance_ratio)
    print("RESULT operator=attention " + " ".join(f"{key}={value}" for key, value in fields.items()))


def save_samples(implementation: Implementation, args: argparse.Namespace, timing: Timing) -> None:
    if args.samples is None:
        return
    fields = metadata(implementation, args)
    with args.samples.open("w", newline="") as file:
        writer = csv.DictWriter(file, fieldnames=(*fields.keys(), "sample", "ms"))
        writer.writeheader()
        for sample, ms in enumerate(timing.samples):
            writer.writerow(dict(fields, sample=sample, ms=ms))


def main() -> int:
    args = parse_args()
    if not torch.cuda.is_available():
        raise RuntimeError("PyTorch cannot see a ROCm GPU")
    print(f'ENV runtime=triton gpu="{torch.cuda.get_device_name(0)}" '
          f"torch={torch.__version__} torch_hip={torch.version.hip} triton={triton.__version__}")
    host_q, host_k, host_v = make_host_input(args.seq, args.dim, args.seed, args.input)
    reference = cpu_reference(host_q, host_k, host_v)
    q, k, v = (tensor.to(device="cuda") for tensor in (host_q, host_k, host_v))
    output = torch.empty_like(q)
    # Keep allocation identical between versions; only materialized kernels
    # access scores/P. This does not demonstrate lower peak device allocation.
    scores = torch.empty((args.seq, args.seq), dtype=torch.float32, device="cuda")
    probabilities = torch.empty_like(scores)
    all_correct = True
    for implementation in selected_implementations(args):
        for tensor in (output, scores, probabilities):
            tensor.fill_(float("nan"))
        launch(implementation, q, k, v, output, scores, probabilities)
        torch.cuda.synchronize()
        precheck = validate(output, reference)
        if not precheck.correct:
            print_result(implementation, args, precheck, None, "FAIL", "NA")
            all_correct = False
            continue
        timing = time_launch(implementation, args, q, k, v, output, scores, probabilities)
        postcheck = validate(output, reference)
        print_result(implementation, args, postcheck, timing, "OK", "OK" if postcheck.correct else "FAIL")
        all_correct = all_correct and postcheck.correct
        if postcheck.correct:
            save_samples(implementation, args, timing)
    return 0 if all_correct else 1


if __name__ == "__main__":
    raise SystemExit(main())

这份文件包含本路线的输入准备、kernel、启动、校验与计时；后续配置共用它，按函数和运行参数定位改动。

完整程序先实现同样的物化路径。三个阶段各启动 S 个 program，一个 program 负责一个 query 行：

阶段	program 内的工作	写出什么
materialized_scores_kernel	以 32 个 key 为一块，计算与当前 query 的点积	当前 query 的整行 Scores
materialized_softmax_kernel	读取一行 Scores，求稳定 Softmax	整行 P
materialized_pv_kernel	以 32 个 key 为一块，合并概率乘 value 的贡献	整行 O

QK 和 PV 都使用 FP32 乘法与 tl.sum，保持本轮精度语义；这里没有通过切换矩阵精度来制造更快的基线。key tile 固定为 32，num_warps=4。BLOCK_D 是覆盖 D 的最小二次幂，超出的分量由 mask 保护。

先运行本路线的物化基线：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
test ! -e "$RUN/n128x64-triton-materialized-k32-p1.csv" && \
python chapter12/attention_triton.py \
  --block-k 32 --num-warps 4 --version materialized \
  --seq 128 --dim 64 --warmup 10 \
  --repeat 50 --seed 20260920 --input dyadic \
  --config-id triton-materialized-k32 --samples "$RUN/n128x64-triton-materialized-k32-p1.csv" --process 1

真实输出：triton-materialized-k32，shape=128x64，第 1 个独立进程

ENV runtime=triton gpu="AMD Radeon RX 9070 XT" torch=2.13.0+rocm10.0.0 torch_hip=7.15.26333 triton=3.8.0
RESULT operator=attention config_id=triton-materialized-k32 implementation=triton-materialized runtime=triton shape=128x64 dtype=float32 block=0 block_k=32 block_d=64 num_warps=4 launches=3 warmup=10 repeat=50 seed=20260920 process=1 scope=gpu-event-full-operator input=dyadic input_precision=fp32 atol=2e-05 rtol=0.0002 intermediate_allocated_bytes=131072 intermediate_used_bytes=131072 seq=128 dim=64 programs_per_query=1 timing=gpu-event reference=fp64-attention-to-fp32 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.03192099928855896 median_ms=0.042000500485301018 mean_ms=0.040016719773411751 max_abs_error=1.4668330550193787e-08 max_tolerance_ratio=0.0007148159511427299

检查正确性、有限值、三个阶段与固定 key tile32，再读完整时间。三进程汇总为 37.940 μs，范围 33.760–42.001 μs；这个范围也保留了较慢的进程。

这里观察到的可改动边界很清楚：相邻阶段通过 Scores/P 交接，而每个 query 的输出不依赖其他 query。可以尝试把这三个阶段在同一 program 内连接起来。

三个阶段之间写出并读回中间矩阵。下一轮保留一行一个 program、相同精度、key tile 32 和 4 warps，改成在一个 program 中直接累计输出。

12.5.5 Triton 改动：合并当前 tile 与历史状态

在线 program 先加载当前 query，随后循环读取 K/V tile。tile 内沿 D 求和得到每个 key 的分数，再沿 key 方向求最大值：

scores = tl.sum(k * q[None, :], axis=1) * scale
scores = tl.where(key_mask, scores, -float("inf"))
tile_max = tl.max(scores, axis=0)
next_max = tl.maximum(running_max, tile_max)
history_scale = tl.exp(running_max - next_max)
probabilities = tl.exp(scores - next_max)

axis=1 合并一个 key 的点积分量，axis=0 合并不同 key 的结果。超出序列末尾的 score 设成负无穷，因此指数贡献为 0。probabilities 在这一段还没有归一化，要等全部 tile 处理完才除以总和。

随后合并向量累计和指数和：

accumulator = accumulator * history_scale + tl.sum(
    probabilities[:, None] * values, axis=0
)
running_sum = running_sum * history_scale + tl.sum(probabilities, axis=0)
running_max = next_max

把手算例改成每次处理两个 key，便能看到尾块的行为。第一块包含 key 0、1；第二块只有 key 2 有效，另一个逻辑位置越界。图 12.6 沿用相同输入，逐步核对尾部 mask 和历史重缩放；这里的块大小 2 仅用于手算。

TRITON · 每块两个 key尾块中一个有效 key 与一个 mask

1/6

读取 key 0、1 的数据副本原输入保持不变。逻辑块每行对应一个 key，V 的两列对应两个分量；第一块两项都有效。

按 key 方向合并贡献基准为 2，两个指数是 1 和 exp(−1)。沿 key 方向相加，得到 l≈1.3679、a≈[1,0.7358]。

key 2 有效，key 3 越界第二块仅 key 2 真实存在。无效位置不读取 V，score 设为 −∞；两个真实 V 分量均有效。

新最大值 4，历史 l/a 同乘 exp(−2)当前块出现 score=4。旧分母与两个旧累计一起重缩放，新 key 的贡献还没有加入。

有效项指数为 1，无效项为 0当前指数为 [1,0]。只加入 key 2 的 1 和 [3,1]；mask 位置对分母及两个分量的贡献都为 0。

全部块处理完，得到 O最后将两个累计分量分别除以 l，得到与逐 key 在线处理和直接加权和相同的结果。

读取 key 0、1 的数据副本原输入保持不变。逻辑块每行对应一个 key，V 的两列对应两个分量；第一块两项都有效。

步骤 1 / 6第一块

这里 D=2 的两个分量都有效，只演示 key 尾部。逻辑块不表示线程或 lane 的排布。

图 12.6　每行对应一个 key，V 的两列对应两个分量。尾块的无效 score 为 −∞，指数为 0，对分母和两个累计分量都没有贡献。本例两个 V 分量均有效；这些逻辑行列不表示硬件 lane 的排布。

这次融合减少了中间数组读写与 kernel 调用，同时引入跨 tile 的历史状态重缩放，并改变归约的组织。因此物化与在线的时间差体现的是整个设计选择，不能只解释成少读写了多少 Byte。

先用 rising-max 穿过多个 key tile，检查历史状态重缩放：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
python chapter12/attention_triton.py \
  --block-k 32 --num-warps 4 --version online \
  --seq 65 --dim 64 --warmup 0 \
  --repeat 1 --seed 20260920 --input rising-max \
  --config-id triton-online-k32

正确性输出：Triton online，65×64，rising-max

ENV runtime=triton gpu="AMD Radeon RX 9070 XT" torch=2.13.0+rocm10.0.0 torch_hip=7.15.26333 triton=3.8.0
RESULT operator=attention config_id=triton-online-k32 implementation=triton-online runtime=triton shape=65x64 dtype=float32 block=0 block_k=32 block_d=64 num_warps=4 launches=1 warmup=0 repeat=1 seed=20260920 process=1 scope=gpu-event-full-operator input=rising-max input_precision=fp32 atol=2e-05 rtol=0.0002 intermediate_allocated_bytes=33800 intermediate_used_bytes=0 seq=65 dim=64 programs_per_query=1 timing=gpu-event reference=fp64-attention-to-fp32 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.094791002571582794 median_ms=0.094791002571582794 mean_ms=0.094791002571582794 max_abs_error=2.9802322387695312e-08 max_tolerance_ratio=0.0008430309877524796

再检查单 key 的直接答案：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
python chapter12/attention_triton.py \
  --block-k 32 --num-warps 4 --version online \
  --seq 1 --dim 1 --warmup 0 \
  --repeat 1 --seed 20260920 --input dyadic \
  --config-id triton-online-k32

正确性输出：Triton online，单 key

ENV runtime=triton gpu="AMD Radeon RX 9070 XT" torch=2.13.0+rocm10.0.0 torch_hip=7.15.26333 triton=3.8.0
RESULT operator=attention config_id=triton-online-k32 implementation=triton-online runtime=triton shape=1x1 dtype=float32 block=0 block_k=32 block_d=1 num_warps=4 launches=1 warmup=0 repeat=1 seed=20260920 process=1 scope=gpu-event-full-operator input=dyadic input_precision=fp32 atol=2e-05 rtol=0.0002 intermediate_allocated_bytes=8 intermediate_used_bytes=0 seq=1 dim=1 programs_per_query=1 timing=gpu-event reference=fp64-attention-to-fp32 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.095991000533103943 median_ms=0.095991000533103943 mean_ms=0.095991000533103943 max_abs_error=0.0 max_tolerance_ratio=0.0

这里仅读正确性。正式比较回到相同 dyadic 输入，保持 key tile32 与 4 warps：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
test ! -e "$RUN/n128x64-triton-online-k32-p1.csv" && \
python chapter12/attention_triton.py \
  --block-k 32 --num-warps 4 --version online \
  --seq 128 --dim 64 --warmup 10 \
  --repeat 50 --seed 20260920 --input dyadic \
  --config-id triton-online-k32 --samples "$RUN/n128x64-triton-online-k32-p1.csv" --process 1

真实输出：triton-online-k32，shape=128x64，第 1 个独立进程

ENV runtime=triton gpu="AMD Radeon RX 9070 XT" torch=2.13.0+rocm10.0.0 torch_hip=7.15.26333 triton=3.8.0
RESULT operator=attention config_id=triton-online-k32 implementation=triton-online runtime=triton shape=128x64 dtype=float32 block=0 block_k=32 block_d=64 num_warps=4 launches=1 warmup=10 repeat=50 seed=20260920 process=1 scope=gpu-event-full-operator input=dyadic input_precision=fp32 atol=2e-05 rtol=0.0002 intermediate_allocated_bytes=131072 intermediate_used_bytes=0 seq=128 dim=64 programs_per_query=1 timing=gpu-event reference=fp64-attention-to-fp32 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.013919999822974205 median_ms=0.014279999770224094 mean_ms=0.016540240030735731 max_abs_error=1.210719347000122e-08 max_tolerance_ratio=0.0005475996007725953

先核对预检、后检与容差，再比较完整时间。源码可以确认中间数组不再读写；真实性能由以下三进程图判断。

图 12.7　本轮固定 key tile 32、4 warps 和一行一个 program，比较三次 kernel 调用的物化版本与一次调用的在线版本。

在线 tile 32 为 14.720 μs，范围 14.280–14.761 μs，相对本路线物化基线约快 2.58 倍。独立确认中，物化为 33.641 μs、在线 tile 32 为 14.860 μs，主要改进再次出现。因此保留在线方案，再检查每轮 key 数是否值得调整。

12.5.6 Triton 调参：固定分工，改变 key tile

在线公式已经不变，现在只把 BLOCK_K 在 16、32、64 中调整。对 S=128，分别需要 8、4、2 轮。其他参数保持一致：一个 program 负责一个 query，num_warps=4，BLOCK_D=64。

更大的 key tile 减少循环和历史状态更新次数，也扩大每轮同时参与表达式的 K/V 数据。它可能改变编译器的寄存器安排，不能从循环减少就判断时间一定下降。

复用 tile32 的记录，只追加 16/64 两个候选：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
test ! -e "$RUN/n128x64-triton-online-k16-p1.csv" && \
python chapter12/attention_triton.py \
  --block-k 16 --num-warps 4 --version online \
  --seq 128 --dim 64 --warmup 10 \
  --repeat 50 --seed 20260920 --input dyadic \
  --config-id triton-online-k16 --samples "$RUN/n128x64-triton-online-k16-p1.csv" --process 1

真实输出：triton-online-k16，shape=128x64，第 1 个独立进程

ENV runtime=triton gpu="AMD Radeon RX 9070 XT" torch=2.13.0+rocm10.0.0 torch_hip=7.15.26333 triton=3.8.0
RESULT operator=attention config_id=triton-online-k16 implementation=triton-online runtime=triton shape=128x64 dtype=float32 block=0 block_k=16 block_d=64 num_warps=4 launches=1 warmup=10 repeat=50 seed=20260920 process=1 scope=gpu-event-full-operator input=dyadic input_precision=fp32 atol=2e-05 rtol=0.0002 intermediate_allocated_bytes=131072 intermediate_used_bytes=0 seq=128 dim=64 programs_per_query=1 timing=gpu-event reference=fp64-attention-to-fp32 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.017279999330639839 median_ms=0.018039999529719353 mean_ms=0.01874909982085228 max_abs_error=1.1059455573558807e-08 max_tolerance_ratio=0.0005446146559292655

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
test ! -e "$RUN/n128x64-triton-online-k64-p1.csv" && \
python chapter12/attention_triton.py \
  --block-k 64 --num-warps 4 --version online \
  --seq 128 --dim 64 --warmup 10 \
  --repeat 50 --seed 20260920 --input dyadic \
  --config-id triton-online-k64 --samples "$RUN/n128x64-triton-online-k64-p1.csv" --process 1

真实输出：triton-online-k64，shape=128x64，第 1 个独立进程

ENV runtime=triton gpu="AMD Radeon RX 9070 XT" torch=2.13.0+rocm10.0.0 torch_hip=7.15.26333 triton=3.8.0
RESULT operator=attention config_id=triton-online-k64 implementation=triton-online runtime=triton shape=128x64 dtype=float32 block=0 block_k=64 block_d=64 num_warps=4 launches=1 warmup=10 repeat=50 seed=20260920 process=1 scope=gpu-event-full-operator input=dyadic input_precision=fp32 atol=2e-05 rtol=0.0002 intermediate_allocated_bytes=131072 intermediate_used_bytes=0 seq=128 dim=64 programs_per_query=1 timing=gpu-event reference=fp64-attention-to-fp32 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.013519999571144581 median_ms=0.014279999770224094 mean_ms=0.015066539905965327 max_abs_error=1.0710209608078003e-08 max_tolerance_ratio=0.0005126649854666966

检查 block_k 变化，同时确认 shape、warps、输入和精度保持不变；各项均通过正确性后，再看独立进程范围。

图 12.8　保留物化原始基线，在线三项只改变 key tile。两个面板采用同一时间刻度；tile 32 是上一轮采用的版本。

tile 16、32、64 分别为 18.040、14.720、14.280 μs。从 16 到 32 的差异清楚，但 64 相对当前 32 只低约 3%，两组进程范围重叠。独立确认中，32 为 14.860 μs（14.760–15.900），64 为 14.300 μs（14.281–15.500），范围仍重叠。本章主形状保留 tile 32，把 64 留作候选，不把更小的中位数直接写成稳定优势。

这里没有新的证据要求为每档 tile 再采一次 profiler。先保留接近候选并复测；若之后出现明确异常，再针对循环、资源或提交方式提出单独问题。

再把 S 改成不能整除 tile 的长度。最后一轮只有一部分 key 有效，key_mask 必须同时保护 K/V 加载，并把无效 score 排除在 Softmax 之外。把无效 score 填 0 会给它们错误地分配指数权重。维度尾部也必须单独用 d < DIM 保护；一个 mask 不能代替另一个。

可以用小输入单独验证两种尾部。仍仅读正确性，不把单次检查时间放入排名：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
python chapter12/attention_triton.py \
  --block-k 32 --num-warps 4 --version online \
  --seq 17 --dim 31 --warmup 0 \
  --repeat 1 --seed 20260920 --input dyadic \
  --config-id triton-online-k32

正确性输出：Triton，17 个 key 与 31 个分量

ENV runtime=triton gpu="AMD Radeon RX 9070 XT" torch=2.13.0+rocm10.0.0 torch_hip=7.15.26333 triton=3.8.0
RESULT operator=attention config_id=triton-online-k32 implementation=triton-online runtime=triton shape=17x31 dtype=float32 block=0 block_k=32 block_d=32 num_warps=4 launches=1 warmup=0 repeat=1 seed=20260920 process=1 scope=gpu-event-full-operator input=dyadic input_precision=fp32 atol=2e-05 rtol=0.0002 intermediate_allocated_bytes=2312 intermediate_used_bytes=0 seq=17 dim=31 programs_per_query=1 timing=gpu-event reference=fp64-attention-to-fp32 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.095440998673439026 median_ms=0.095440998673439026 mean_ms=0.095440998673439026 max_abs_error=2.9802322387695312e-08 max_tolerance_ratio=0.000989819581660001

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
python chapter12/attention_triton.py \
  --block-k 32 --num-warps 4 --version online \
  --seq 33 --dim 65 --warmup 0 \
  --repeat 1 --seed 20260920 --input dyadic \
  --config-id triton-online-k32

正确性输出：Triton，33 个 key 与 65 个分量

ENV runtime=triton gpu="AMD Radeon RX 9070 XT" torch=2.13.0+rocm10.0.0 torch_hip=7.15.26333 triton=3.8.0
RESULT operator=attention config_id=triton-online-k32 implementation=triton-online runtime=triton shape=33x65 dtype=float32 block=0 block_k=32 block_d=128 num_warps=4 launches=1 warmup=0 repeat=1 seed=20260920 process=1 scope=gpu-event-full-operator input=dyadic input_precision=fp32 atol=2e-05 rtol=0.0002 intermediate_allocated_bytes=8712 intermediate_used_bytes=0 seq=33 dim=65 programs_per_query=1 timing=gpu-event reference=fp64-attention-to-fp32 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.092570997774600983 median_ms=0.092570997774600983 mean_ms=0.092570997774600983 max_abs_error=1.862645149230957e-08 max_tolerance_ratio=0.000695996040016387

先核对检查通过，再看 block_d 随 D 补齐；key 与分量两个维度都必须保护。

另外两组性能输入为 S=256,D=64 与 S=128,D=65，每组仍运行全部六配置、三个独立进程。前者的在线 tile 32/64 为 20.800/19.160 μs，后者则为 24.061/25.440 μs，次序改变了。

D 从 64 到 65 时，覆盖它的 BLOCK_D 从 64 扩大到 128，逻辑块中无效的分量明显增加；这是下一步值得研究的变化，不能仅凭现有时间给它单独分摊成本。两组数据只用于观察形状迁移，尚未分别做独立候选确认，完整比较见 图 12.10。

12.6 完整路径比较与方案选择

物化与在线计算保留相同的 Attention 语义，却用不同方式安排中间结果。一次融合是否值得，要同时检查省掉了什么、增加了什么，并用同范围的计时作选择。

路线与主形状 S×D=128×64	主扫描完整时间（μs）	本轮决定
HIP materialized / block256	59.681	保留为 HIP 方案
HIP online / block256	230.643	保留正确实现与负结果，继续诊断执行安排
Triton materialized / key32	37.940	保留为路线原始基线
Triton online / key32	14.720	采用在线版本，保留当前 tile32
Triton online / key64	14.280	候选与 tile32 的范围重叠，暂不替换默认选择

图 12.9　RX 9070 XT、ROCm 10.0，S=128、D=64，完整路径 event。两条路线分开比较；每项为三个新进程，不与主扫描合并。

独立确认保留了所有六个配置，包括较慢的在线 HIP 和接近的 Triton tile32/64。HIP 的负结果、Triton 在线相对物化的收益都复现了；tile32 与 tile64 的差异仍不足以据此作稳定选择。

图 12.10　S×D=256×64 与 128×65 各测全部六项。每个面板单独标注路线与形状，刻度独立，不能跨面板直接比较柱长。

两组附加形状中，HIP 物化仍快于当前在线实现，Triton 在线仍快于本路线物化实现；哪种在线 key tile 更合适则发生变化。形状迁移保留了全部配置，未只展示有利的一组。

本章的选择范围是已经检查过的 shape、输入、精度与设备。它不等于生产 Attention 的普遍结论：当前实现没有 query tile 之间的高级复用，没有混合精度矩阵路径，也没有 batch/head/causal。换成其他输入范围、序列长度或融合边界后，需要重新选择基线并验证。

12.7 运行与诊断

前面的命令便于逐项理解。要完成三个独立进程的整组比较，仍在已激活的 code/part2-kernels/ 中，使用批量入口；它以尚不存在的 $RUN/rounds 子目录保存本轮结果：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
python chapter12/run_rounds.py --output "$RUN/rounds" --phase main && \
python chapter12/plot_rounds.py \
  --evidence "$RUN/rounds/summary" --out-dir "$RUN/figures/main" \
  --round hip-fusion --round triton-fusion --round triton-tile

主阶段固定 S=128,D=64，先完成 36 项边界与 rising-max 检查，再运行六配置。 每项三个独立进程，各预热 10 次、计时 50 次，轮换候选顺序；先检查正确性，再汇总性能。main 不采集 profiler。图中心为三个进程中位数的中位数，误差线为它们的最小值到最大值，不是置信区间。

需要检查形状迁移时，再运行 validation：

python chapter12/run_rounds.py --output "$RUN/rounds" --phase validation && \
python chapter12/plot_rounds.py \
  --evidence "$RUN/rounds/summary" --out-dir "$RUN/figures/validation" \
  --round shapes

两组附加形状是 S×D=256×64 与 128×65，每组测全部六项。这一步继续使用本轮冻结的源码和配置，并检查实际环境。两条绘图命令都显式读取你的本轮 summary，不会用教程发布的成绩替缺失结果补值。

选出候选后，再以独立目录确认：

python chapter12/verify_rounds.py \
  --frozen-run "$RUN/rounds" --output "$RUN/confirmation"

确认保留 HIP 两版、Triton 物化与 online16，再加最快在线配置及距它不超过 5% 的候选，重复项只运行一次。本章测试因此确认了全部六项配置。确认数据保存在自己的 summary/summary.csv、summary/process-summary.csv 与 confirmation.json 中，不自动并入父实验。先比较同形状的新旧进程范围，再决定是否保留候选；不要复制教程中的确认文件到新目录拼图。

程序把逐次计时写入 samples/，把检查与终端记录写入 logs/；source/ 保存本轮源码，commands.jsonl 保存实际参数。已有主目录拒绝被新的 main 覆盖；修改代码、改变条件或重新采一轮时，新建 RUN。第一次编译和 Triton JIT 在稳态计时之外，后续终端继续使用本篇环境。

本章图使用 chapter12/evidence/rounds/ 的三进程实测与独立确认；折叠终端输出来自其中第 1 个进程的原始日志，保存在 chapter12/evidence/walkthrough/。一次进程的中位数不必等于图的中心。按需 trace 的采集命令与阅读报告已经放在对应问题旁；批量 --phase profile 仅供需要完整诊断矩阵时选用。

本章实测条件为 RX 9070 XT、ROCm SDK 10.0.0、原生 Ubuntu 24.04.5，HIP 7.15.26333、Triton 3.8.0；GPU 非独占、不锁频，输入数组重复使用且不清缓存。PMC、ATT 等工具并非本章操作前提，当前可用能力与限制见第 8 章。

12.8 练习与下一轮实验

1. 仍用分数 [2,1,4] 和三个 value，手算第三轮的 $\alpha ,l,a$。如果分母重缩放而分子没有，哪些输出分量会改变？
2. 把 key 顺序改成 [4,2,1]，同时重排对应 value。数学结果应怎样变化？FP32 的最后几位是否必须一致？
3. 对 S=33，分别画出 key tile 16 与 32 的有效位置。保持其他参数不变，先验证答案，再记录完整时间，并说明变化是否超过独立进程间的范围。
4. HIP 选读：按 图 12.4 标出每个屏障保护的读写依赖。若下一轮尝试一次处理多个 key，哪些状态可以局部合并，哪些依赖仍必须等待？先说明设计，再改代码。

本章小结

Attention 的输出依赖所有 key 的权重，但不一定需要保存全部分数。在线算法通过最大值、指数和及向量累计保留历史贡献；最大值改变时，分母和分子一起重缩放。

融合改变的不只有中间写回，也可能改变并行分工、同步次数和同时保存的数据量。HIP 的物化与在线对照、Triton 的融合与 key tile 扫描都要由完整计时和正确性检查决定是否采用。下一章会把这套方法用于更紧凑的 RMSNorm，让你独立完成一次候选配置的选择与复测。

延伸阅读

• FlashAttention 论文：理解分块和减少中间读写的设计。本章只保留便于教学观察的 FP32 前向版本，不复现完整论文 kernel。
• Triton Fused Attention 教程：继续阅读更完整的分块实现；其中部分配置与硬件路径有额外前提，不能直接当作本章 gfx1201 的已验证选项。
• 附录 D：HIP 与 Triton 的编程范式：回顾 block 协作与 program 的逻辑数据块。