第10章 10Normalization：归一化算子

本章导读

第 9 章把多个输入合成一个和。现在希望把一行分数变成一行非负权重，让它们加起来等于 1。Softmax 先取指数，再除以整行指数的和；因此每个输出既需要自己的输入，也需要整行共同算出的分母。

本章用 [1000, 1001, 1002] 看清稳定公式，再分别研究两个问题：一行里的计算怎样分工，中间结果是否值得留在一次 kernel 内完成？ HIP 先建立每行一个 block 的三阶段基线，再比较融合；Triton 从一行一个 program 出发，分别比较参与的 warps 数和逻辑向量宽度。

配套代码 编程范式附录

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

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

• 先理解算法：阅读行级语义、数值稳定性与中间数据流。
• 选择 HIP：进入 HIP 路线，先确认行内协作，再验证融合是否划算。
• 选择 Triton：进入 Triton 路线，用固定参数的对照逐步选择执行配置。

配套代码在 code/part2-kernels/chapter10/，环境沿用 Radeon RX 9070 XT（gfx1201）+ ROCm 10.0 + 原生 Ubuntu 24.04。归约、mask 与 GPU event 分别沿用第 9、8、5 章；动画展示手算步骤，不表示真实执行速度。

10.1 从一行分数到一行概率

输入一行 [2, 1, 0]，先逐位置取自然指数，再除以这些指数的总和：

输入 x	exp(x)，约值	除以总和 11.1074 后，约值
2	7.3891	0.6652
1	2.7183	0.2447
0	1.0000	0.0900

分数越大，得到的权重越大。未舍入输出的和为 1，因此可以解释成一组概率；表中四位小数之和为 0.9999，只是显示舍入。

对 R×C 的矩阵 X，本章沿每一行的列维度计算：

$$Y_{r,c}=\frac{e^{X_{r,c}}}{\sum _{j=0}^{C-1}{e}^{X_{r,j}}},0\le r<R,0\le c<C.$$

r 是行号，c 是列号。一个输出的分子来自当前位置，分母却来自整行；同一行内需要协作，不同行互不依赖。我们不会为整个矩阵只计算一个共同分母。

10.2 平移输入，避免指数溢出

把输入换成 [1000, 1001, 1002]。直接计算 exp(1000) 已远超 FP32 的有限范围，最后可能得到 inf/inf。反过来，[-1002, -1001, -1000] 的指数可能全部下溢为 0，分母也变成 0。

整行减去同一个常数，可以在数学上保留原来的比例：

$$\frac{e^{x_c-a}}{\sum _j{e}^{x_j-a}}=\frac{e^{-a}{e}^{x_c}}{e^{-a}\sum _j{e}^{x_j}}=\frac{e^{x_c}}{\sum _j{e}^{x_j}}.$$

选择 a=max(x)，最大的指数自变量就变成 0：

步骤	[1000,1001,1002] 的结果
最大值 m	1002
平移 x−m	[-2, -1, 0]
指数 p=exp(x−m)	约 [0.1353, 0.3679, 1.0000]
分母 s=sum(p)	约 1.5032
输出 y=p/s	约 [0.0900, 0.2447, 0.6652]

手算示例 · 一行三个值整行归约，再将标量分发给各位置

1/6

跟踪同一行的三个数输入是 1000、1001、1002。原值一直保留在顶行；最大值、分母与输出尚未产生。

三个输入共同决定一个最大值箭头把整行输入送到最大值归约，完成后得到 m=1002。它属于整行，下一步每个位置都需要这个数。

每项输入减去同一个 m把 1002 的副本分发给各位置，分别计算 1000−1002、1001−1002、1002−1002，得到 −2、−1、0。顶行原值保留。

各位置得到稳定指数对 −2、−1、0 分别取指数，得到约 0.1353、0.3679、1.0000。此时还没有分母，也不能写出最终概率。

三个指数共同决定一个分母把整行指数相加，得到 s≈1.5032；原来的指数仍保留，用于下一步的除法。

每项指数除以同一个 s把分母分发给各位置，计算 p/s，得到约 0.0900、0.2447、0.6652。未舍入值的和为 1，四位小数的和为 0.9999，差别来自显示舍入。

跟踪同一行的三个数输入是 1000、1001、1002。原值一直保留在顶行；最大值、分母与输出尚未产生。

步骤 1 / 6输入

箭头表示数据依赖与值的副本，不指定物理线程或存储位置；动画速度不表示 GPU 耗时。

图 10.1　同一行先汇入最大值 m，再把 m 分发给各位置；指数求和得到 s 后，所有位置再使用同一个分母。数字在运算完成时更新，图中小数保留四位。

图 10.1 中，先停在“汇入 max”，确认三个输入共同产生一个 1002；再进入“分发 m”，看同一个值怎样用于三次减法。分母也遵循相同的汇入与分发关系。最大值和分母各只有一个，却要供整行共同使用。这正是两次归约的职责。减去最大值后，至少一项指数为 1；其他项处在 0 与 1 之间，因此不会因为全部指数下溢而得到零分母。

所有版本采用相同的稳定形式：

$$m_r=\underset{j}{max}{X}_{r,j},p_{r,c}=e^{X_{r,c}-m_r},s_r=\sum _j{p}_{r,j},Y_{r,c}=p_{r,c}/s_r.$$

本章输入限定为有限 FP32 值。减去最大值解决了直接指数的主要范围问题，并不消除舍入误差；非常小的概率仍可能下溢为 0。含 NaN 或无穷的输入需要另行规定语义。

10.3 中间结果与计算边界

稳定公式串起了四个逻辑操作：

操作	已学过的模式	下一步需要什么
max(x)	Reduction	整行共同使用的最大值
exp(x−m)	Element-Wise	每列一个稳定指数
sum(p)	Reduction	整行共同使用的分母
p/s	Element-Wise	每列一个输出

四个逻辑操作不要求启动四个 kernel。一个直接方案用三个阶段：求最大值；计算并保存指数，同时求和；读取指数与行和，完成归一化。后一个阶段要读前一个阶段的结果，所以它们通过全局数组连接。

图 10.2　上方用中间数组连接三次调用；下方尝试在一次调用中完成。融合版如何保存或重算指数，要到具体实现中判断。

图 10.2 给出一个优化方向：少一次中间数组往返，是否值得？主输入为 4096×1024，单个 FP32 指数数组 exp_tmp 就有 16 MiB，完整写入再读取对应 32 MiB 的逻辑访问。

但省掉这个数组后，指数仍要用于最后的除法。可以保留它，也可以重新计算。HIP 本章采用重算，Triton 在 program 的逻辑表达式中复用指数值。我们会分别检查各自新增和减少的工作，再用完整计时作选择。

10.4 正确性参考与计时范围

两条路线使用同一输入语义：连续的 R×C FP32 矩阵，沿最后一维做 Softmax。测试值随行列位置变化，并按行叠加 +1000、−1000、0，以覆盖直接取指数容易出问题的范围；本轮固定 seed=20260920。

检查项目	本章约定
输入与输出	有限 FP32 输入，同形状 FP32 输出
正确性参考	CPU 按 FP64 稳定公式计算，再转为 FP32 与输出比较
逐元素误差	输出必须有限，最大绝对误差不超过 2e-5
行和误差	每行与 1 的差不超过 2e-5；两条路线均用 FP64 累加输出行和
检查时机	计时前检查完整输出；计时后检查最后一次完整输出
初始化检查	首次输出预填 NaN，使未写入位置能被发现

正确性参考负责判断答案；HIP 的路线基线是 cooperative 三阶段，Triton 的路线基线是紧凑逻辑宽度与 4 warps。CPU reference 不作为 GPU 性能基线。

本章基准测试的主形状定为 4096×1024。为了检验不同尺寸下的泛化表现，我们还会测试 4096×1023（比主形状少一列的尾部边界）、4096×1025（超过 1024 列后需要额外补齐与掩码保护的边界）以及行数较少的 128×1024；主矩阵与这些补充形状共涵盖 75 个独立进程、3,750 个计时样本。此外，在微小尺寸（如 1×1、3×33、3×257、3×1025）上验证全部七种配置的数值正确性（共 28 项检查通过）。这里验证的是带较大正负偏移的代表性有限输入，并不意味着任意极端数值输入都能获得绝对数学保证。

完整 Softmax 的 event 区间从第一个 kernel 之前开始，到最后一个 kernel 之后结束：三阶段包含全部三次调用，融合包含一次调用。两条路线每轮都等待结束 event 完成，再读取该轮时间。内存分配、初始数据拷贝、CPU reference、输出检查和 Triton 首次 JIT 均在稳态计时之外。

正式采集于 2026-09-20，Radeon RX 9070 XT（gfx1201）、ROCm SDK 10.0.0、原生 Ubuntu 24.04，HIP 为 7.15.26333，Triton 为 3.8.0。每配置三个独立进程，各预热 10 次、计时 50 次，轮换执行顺序；同一进程复用数组、不主动刷缓存，设备保留桌面与已有上下文。短 event 区间仍可能包括 CPU 提交空隙。

正式计时不打开 profiler，trace 另行采集。结果图以各独立进程 median 的中位数作为中心，范围表示这些 median 的最小值到最大值，不是置信区间。先检查完整时间，再用阶段诊断解释下一步要研究什么。

10.5 HIP 与 Triton 优化路线

两条路线分别建立自然的基线，再围绕一个问题修改。先选一种语言读完整即可。

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

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

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

mkdir -p chapter10/results &&
RUN=$(mktemp -d chapter10/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。

本章把三个问题分开：答案与行和是否正确，完整 Softmax 是否更快，以及融合后程序是否保持了计划中的行分工。前两个问题每轮都检查，第三个问题只在融合对照中用一组 trace 核对。

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

10.5.1 公式入口：展开三个阶段

完整源码：code/part2-kernels/chapter10/softmax_hip.hip

// Chapter 10: numerically stable row-wise Softmax in HIP.
//
// The baseline deliberately exposes three global-memory stages. The fused
// version assigns one block to one row and keeps both reductions in LDS.

#include <hip/hip_runtime.h>

#include <algorithm>
#include <cmath>
#include <cstddef>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <exception>
#include <fstream>
#include <iomanip>
#include <limits>
#include <numeric>
#include <string>
#include <vector>

#define HIP_CHECK(call)                                                        \
    do {                                                                       \
        const hipError_t error = (call);                                       \
        if (error != hipSuccess) {                                             \
            std::fprintf(stderr, "HIP error at %s:%d: %s (%s)\n",             \
                         __FILE__, __LINE__, hipGetErrorString(error), #call);  \
            std::exit(EXIT_FAILURE);                                           \
        }                                                                      \
    } while (0)

namespace {

constexpr float kDeviceNegativeInfinity = -std::numeric_limits<float>::infinity();
constexpr double kAbsoluteTolerance = 2.0e-5;
constexpr double kRowSumTolerance = 2.0e-5;

__global__ void row_max_serial(const float* __restrict__ input,
                               float* __restrict__ row_max,
                               std::size_t rows,
                               std::size_t columns) {
    const std::size_t row =
        static_cast<std::size_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    if (row >= rows) {
        return;
    }

    float maximum = kDeviceNegativeInfinity;
    const std::size_t base = row * columns;
    for (std::size_t column = 0; column < columns; ++column) {
        maximum = fmaxf(maximum, input[base + column]);
    }
    row_max[row] = maximum;
}

__global__ void row_exp_sum_serial(const float* __restrict__ input,
                                   const float* __restrict__ row_max,
                                   float* __restrict__ exponentials,
                                   float* __restrict__ row_sum,
                                   std::size_t rows,
                                   std::size_t columns) {
    const std::size_t row =
        static_cast<std::size_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    if (row >= rows) {
        return;
    }

    const std::size_t base = row * columns;
    float sum = 0.0F;
    for (std::size_t column = 0; column < columns; ++column) {
        const float value = expf(input[base + column] - row_max[row]);
        exponentials[base + column] = value;
        sum += value;
    }
    row_sum[row] = sum;
}

__global__ void normalize_rows(const float* __restrict__ exponentials,
                               const float* __restrict__ row_sum,
                               float* __restrict__ output,
                               std::size_t elements,
                               std::size_t columns) {
    const std::size_t index =
        static_cast<std::size_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    if (index < elements) {
        output[index] = exponentials[index] / row_sum[index / columns];
    }
}

// Bridge experiment: preserve three stages while assigning one block per row.
// The max/sum trees and column-stride mapping match softmax_fused_lds.
__global__ void row_max_block(const float* __restrict__ input,
                             float* __restrict__ row_max,
                             std::size_t columns) {
    extern __shared__ float shared[];
    const unsigned int lane = threadIdx.x;
    const std::size_t base = static_cast<std::size_t>(blockIdx.x) * columns;
    float local = kDeviceNegativeInfinity;
    for (std::size_t column = lane; column < columns; column += blockDim.x)
        local = fmaxf(local, input[base + column]);
    shared[lane] = local;
    __syncthreads();
    for (unsigned int stride = blockDim.x / 2; stride > 0; stride /= 2) {
        if (lane < stride) shared[lane] = fmaxf(shared[lane], shared[lane + stride]);
        __syncthreads();
    }
    if (lane == 0) row_max[blockIdx.x] = shared[0];
}

__global__ void row_exp_sum_block(const float* __restrict__ input,
                                 const float* __restrict__ row_max,
                                 float* __restrict__ exponentials,
                                 float* __restrict__ row_sum,
                                 std::size_t columns) {
    extern __shared__ float shared[];
    const unsigned int lane = threadIdx.x;
    const std::size_t base = static_cast<std::size_t>(blockIdx.x) * columns;
    const float maximum = row_max[blockIdx.x];
    float local = 0.0F;
    for (std::size_t column = lane; column < columns; column += blockDim.x) {
        const float value = expf(input[base + column] - maximum);
        exponentials[base + column] = value;
        local += value;
    }
    shared[lane] = local;
    __syncthreads();
    for (unsigned int stride = blockDim.x / 2; stride > 0; stride /= 2) {
        if (lane < stride) shared[lane] += shared[lane + stride];
        __syncthreads();
    }
    if (lane == 0) row_sum[blockIdx.x] = shared[0];
}

__global__ void normalize_block(const float* __restrict__ exponentials,
                               const float* __restrict__ row_sum,
                               float* __restrict__ output,
                               std::size_t columns) {
    const std::size_t base = static_cast<std::size_t>(blockIdx.x) * columns;
    const float denominator = row_sum[blockIdx.x];
    for (std::size_t column = threadIdx.x; column < columns; column += blockDim.x)
        output[base + column] = exponentials[base + column] / denominator;
}

__global__ void softmax_fused_lds(const float* __restrict__ input,
                                  float* __restrict__ output,
                                  std::size_t rows,
                                  std::size_t columns) {
    const std::size_t row = blockIdx.x;
    if (row >= rows) {
        return;
    }

    extern __shared__ float shared[];
    const unsigned int lane = threadIdx.x;
    const std::size_t base = row * columns;

    float local_maximum = kDeviceNegativeInfinity;
    for (std::size_t column = lane; column < columns; column += blockDim.x) {
        local_maximum = fmaxf(local_maximum, input[base + column]);
    }
    shared[lane] = local_maximum;
    __syncthreads();

    for (unsigned int stride = blockDim.x / 2; stride > 0; stride /= 2) {
        if (lane < stride) {
            shared[lane] = fmaxf(shared[lane], shared[lane + stride]);
        }
        __syncthreads();
    }
    const float maximum = shared[0];
    // Every thread must read the maximum before LDS is reused for the sum.
    __syncthreads();

    float local_sum = 0.0F;
    for (std::size_t column = lane; column < columns; column += blockDim.x) {
        local_sum += expf(input[base + column] - maximum);
    }
    shared[lane] = local_sum;
    __syncthreads();

    for (unsigned int stride = blockDim.x / 2; stride > 0; stride /= 2) {
        if (lane < stride) {
            shared[lane] += shared[lane + stride];
        }
        __syncthreads();
    }
    const float denominator = shared[0];

    for (std::size_t column = lane; column < columns; column += blockDim.x) {
        output[base + column] =
            expf(input[base + column] - maximum) / denominator;
    }
}

struct Args {
    std::string version = "all";
    std::size_t rows = 4096;
    std::size_t columns = 1024;
    int block = 256;
    int warmup = 10;
    int repeat = 50;
    std::uint64_t seed = 20260920;
    std::string samples;
    std::string config_id;
    int process = 1;
};

enum class Kind { baseline, cooperative, fused };

struct Implementation {
    const char* name;
    Kind kind;
    int launches;
};

struct Validation {
    bool correct;
    bool finite;
    double max_absolute_error;
    double max_row_sum_error;
    std::size_t first_bad_index;
};

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

[[noreturn]] void usage_error(const std::string& message) {
    std::fprintf(
        stderr,
        "%s\nUsage: softmax_hip [--version all|baseline|cooperative|fused] "
        "[--rows N] [--cols N] [--block N] [--warmup N] "
        "[--repeat N] [--seed N] [--samples CSV] [--config-id ID] [--process N]\n",
        message.c_str());
    std::exit(EXIT_FAILURE);
}

Args parse_args(int argc, char** argv) {
    Args args;
    for (int index = 1; index < argc; ++index) {
        const std::string option = argv[index];
        auto next = [&]() -> std::string {
            if (index + 1 >= argc) {
                usage_error("missing value for " + option);
            }
            return argv[++index];
        };

        try {
            if (option == "--version") {
                args.version = next();
            } else if (option == "--rows") {
                args.rows = std::stoull(next());
            } else if (option == "--cols") {
                args.columns = std::stoull(next());
            } else if (option == "--block") {
                args.block = std::stoi(next());
            } else if (option == "--warmup") {
                args.warmup = std::stoi(next());
            } else if (option == "--repeat") {
                args.repeat = std::stoi(next());
            } else if (option == "--seed") {
                args.seed = std::stoull(next());
            } else if (option == "--samples") {
                args.samples = next();
            } else if (option == "--config-id") {
                args.config_id = next();
            } else if (option == "--process") {
                args.process = std::stoi(next());
            } else {
                usage_error("unknown option: " + option);
            }
        } catch (const std::exception&) {
            usage_error("invalid value for " + option);
        }
    }

    if (args.rows == 0 || args.columns == 0) {
        usage_error("--rows and --cols must be positive");
    }
    if (args.rows > std::numeric_limits<unsigned int>::max()) {
        usage_error("--rows is too large for the fused grid");
    }
    if (args.columns > std::numeric_limits<std::size_t>::max() / args.rows) {
        usage_error("rows * columns overflows size_t");
    }
    if (args.rows * args.columns >
        std::numeric_limits<std::size_t>::max() / sizeof(float)) {
        usage_error("the FP32 tensor byte size overflows size_t");
    }
    if (args.block <= 0 || args.block > 1024 ||
        (args.block & (args.block - 1)) != 0) {
        usage_error("--block must be a power of two between 1 and 1024");
    }
    if (args.warmup < 0 || args.repeat <= 0) {
        usage_error("--warmup must be non-negative and --repeat positive");
    }
    if (args.version != "all" && args.version != "baseline" &&
        args.version != "cooperative" && args.version != "fused") {
        usage_error("unknown version: " + args.version);
    }
    if (!args.samples.empty() && (args.version == "all" || args.config_id.empty()))
        usage_error("--samples requires one version and a --config-id");
    if (args.seed > UINT32_MAX) usage_error("seed must fit uint32");
    if (args.process <= 0 || args.config_id.find_first_not_of("abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789-_.") != std::string::npos)
        usage_error("invalid process or config-id");
    return args;
}

std::vector<Implementation> selected_implementations(
    const std::string& version) {
    const Implementation baseline{"hip-baseline-3kernel", Kind::baseline, 3};
    const Implementation cooperative{"hip-cooperative-3kernel", Kind::cooperative, 3};
    const Implementation fused{"hip-fused-block-lds", Kind::fused, 1};
    if (version == "all") {
        return {baseline, cooperative, fused};
    }
    if (version == "cooperative") return {cooperative};
    return version == "baseline" ? std::vector<Implementation>{baseline}
                                 : std::vector<Implementation>{fused};
}

std::size_t ceil_div(std::size_t numerator, std::size_t denominator) {
    return numerator / denominator + (numerator % denominator != 0);
}

unsigned int checked_grid(std::size_t blocks) {
    if (blocks == 0 || blocks > std::numeric_limits<unsigned int>::max()) {
        usage_error("calculated grid is outside the supported range");
    }
    return static_cast<unsigned int>(blocks);
}

void launch(const Implementation& implementation,
            const Args& args,
            const float* input,
            float* output,
            float* exponentials,
            float* row_max,
            float* row_sum) {
    const dim3 block(static_cast<unsigned int>(args.block));
    const std::size_t elements = args.rows * args.columns;

    if (implementation.kind == Kind::baseline) {
        const dim3 row_grid(checked_grid(ceil_div(args.rows, args.block)));
        const dim3 element_grid(checked_grid(ceil_div(elements, args.block)));
        hipLaunchKernelGGL(row_max_serial, row_grid, block, 0, 0,
                           input, row_max, args.rows, args.columns);
        hipLaunchKernelGGL(row_exp_sum_serial, row_grid, block, 0, 0,
                           input, row_max, exponentials, row_sum,
                           args.rows, args.columns);
        hipLaunchKernelGGL(normalize_rows, element_grid, block, 0, 0,
                           exponentials, row_sum, output,
                           elements, args.columns);
    } else if (implementation.kind == Kind::cooperative) {
        const dim3 row_grid(static_cast<unsigned int>(args.rows));
        const std::size_t lds_bytes = static_cast<std::size_t>(args.block) * sizeof(float);
        hipLaunchKernelGGL(row_max_block, row_grid, block, lds_bytes, 0,
                           input, row_max, args.columns);
        hipLaunchKernelGGL(row_exp_sum_block, row_grid, block, lds_bytes, 0,
                           input, row_max, exponentials, row_sum, args.columns);
        hipLaunchKernelGGL(normalize_block, row_grid, block, 0, 0,
                           exponentials, row_sum, output, args.columns);
    } else {
        const dim3 row_grid(static_cast<unsigned int>(args.rows));
        const std::size_t lds_bytes =
            static_cast<std::size_t>(args.block) * sizeof(float);
        hipLaunchKernelGGL(softmax_fused_lds, row_grid, block, lds_bytes, 0,
                           input, output, args.rows, args.columns);
    }
    HIP_CHECK(hipGetLastError());
}

std::vector<float> make_input(const Args& args) {
    std::vector<float> input(args.rows * args.columns);
    for (std::size_t row = 0; row < args.rows; ++row) {
        const float shift = row % 3 == 0 ? 1000.0F
                           : row % 3 == 1 ? -1000.0F
                                          : 0.0F;
        for (std::size_t column = 0; column < args.columns; ++column) {
            const std::uint64_t mixed =
                static_cast<std::uint64_t>(column) * 17ULL +
                static_cast<std::uint64_t>(row) * 131ULL + args.seed;
            const int centered = static_cast<int>(mixed % 4096ULL) - 2048;
            input[row * args.columns + column] =
                shift + static_cast<float>(centered) / 128.0F;
        }
    }
    return input;
}

std::vector<float> cpu_reference(const std::vector<float>& input,
                                 std::size_t rows,
                                 std::size_t columns) {
    std::vector<float> reference(input.size());
    std::vector<double> exponentials(columns);
    for (std::size_t row = 0; row < rows; ++row) {
        const std::size_t base = row * columns;
        double maximum = -std::numeric_limits<double>::infinity();
        for (std::size_t column = 0; column < columns; ++column) {
            maximum = std::max(maximum, double(input[base + column]));
        }

        double sum = 0.0;
        for (std::size_t column = 0; column < columns; ++column) {
            const double value = std::exp(double(input[base + column]) - maximum);
            exponentials[column] = value;
            sum += value;
        }
        for (std::size_t column = 0; column < columns; ++column) {
            reference[base + column] = float(exponentials[column] / sum);
        }
    }
    return reference;
}

Validation validate(const std::vector<float>& output,
                    const std::vector<float>& reference,
                    std::size_t rows,
                    std::size_t columns) {
    Validation result{true, true, 0.0, 0.0, output.size()};
    for (std::size_t row = 0; row < rows; ++row) {
        double row_sum = 0.0;
        for (std::size_t column = 0; column < columns; ++column) {
            const std::size_t index = row * columns + column;
            result.finite = result.finite && std::isfinite(output[index]);
            row_sum += output[index];
            const double error =
                std::abs(static_cast<double>(output[index]) - reference[index]);
            result.max_absolute_error =
                std::max(result.max_absolute_error, error);
            if ((!std::isfinite(output[index]) ||
                 error > kAbsoluteTolerance) &&
                result.first_bad_index == output.size()) {
                result.first_bad_index = index;
                result.correct = false;
            }
        }
        result.max_row_sum_error =
            std::max(result.max_row_sum_error, std::abs(row_sum - 1.0));
    }
    result.correct = result.correct &&
                     result.max_row_sum_error <= kRowSumTolerance;
    return result;
}

Timing benchmark(const Implementation& implementation,
                 const Args& args,
                 const float* input,
                 float* output,
                 float* exponentials,
                 float* row_max,
                 float* row_sum) {
    for (int iteration = 0; iteration < args.warmup; ++iteration) {
        launch(implementation, args, input, output,
               exponentials, row_max, row_sum);
    }
    HIP_CHECK(hipDeviceSynchronize());

    hipEvent_t start;
    hipEvent_t end;
    HIP_CHECK(hipEventCreate(&start));
    HIP_CHECK(hipEventCreate(&end));
    std::vector<double> times_ms;
    times_ms.reserve(args.repeat);
    for (int iteration = 0; iteration < args.repeat; ++iteration) {
        HIP_CHECK(hipEventRecord(start, nullptr));
        launch(implementation, args, input, output,
               exponentials, row_max, row_sum);
        HIP_CHECK(hipEventRecord(end, nullptr));
        HIP_CHECK(hipEventSynchronize(end));
        float elapsed_ms = 0.0F;
        HIP_CHECK(hipEventElapsedTime(&elapsed_ms, start, end));
        times_ms.push_back(elapsed_ms);
    }
    HIP_CHECK(hipEventDestroy(start));
    HIP_CHECK(hipEventDestroy(end));

    const std::vector<double> samples = times_ms;
    std::sort(times_ms.begin(), times_ms.end());
    const double median = times_ms.size() % 2 == 0
        ? (times_ms[times_ms.size() / 2 - 1] +
           times_ms[times_ms.size() / 2]) / 2.0
        : times_ms[times_ms.size() / 2];
    return Timing{
        times_ms.front(),
        median,
        std::accumulate(times_ms.begin(), times_ms.end(), 0.0) /
            times_ms.size(),
        samples,
    };
}

void print_result(const Implementation& implementation,
                  const Args& args,
                  const Timing* timing,
                  const Validation& validation,
                  const char* precheck,
                  const char* postcheck) {
    const char* status = validation.correct ? "OK" : "FAIL";
    std::printf(
        "RESULT operator=softmax implementation=%s runtime=hip "
        "shape=%zux%zu dtype=float32 block=%d launches=%d lds_bytes=%zu "
        "warmup=%d repeat=%d seed=%llu timing=gpu-event "
        "scope=gpu-event-full-operator input=shifted-dyadic reference=fp64-softmax-to-fp32 "
        "timed=%d correct=%s precheck=%s postcheck=%s ",
        implementation.name, args.rows, args.columns, args.block,
        implementation.launches,
        implementation.kind != Kind::baseline
            ? static_cast<std::size_t>(args.block) * sizeof(float)
            : 0,
        args.warmup, args.repeat,
        static_cast<unsigned long long>(args.seed), timing != nullptr,
        status, precheck, postcheck);
    std::printf("config_id=%s num_warps=0 output_finite=%s ", args.config_id.c_str(), validation.finite ? "OK" : "FAIL");
    if (timing == nullptr) {
        std::printf("min_ms=NA median_ms=NA mean_ms=NA ");
    } else {
        std::printf("min_ms=%.9f median_ms=%.9f mean_ms=%.9f ",
                    timing->minimum_ms, timing->median_ms, timing->mean_ms);
    }
    std::printf("max_abs_error=%.9g max_row_sum_error=%.9g\n",
                validation.max_absolute_error,
                validation.max_row_sum_error);
}

void save_samples(const Implementation& implementation, const Args& args,
                  const Timing& timing) {
    if (args.samples.empty()) return;
    std::ofstream file(args.samples);
    if (!file) usage_error("cannot open sample file");
    file << "config_id,implementation,runtime,shape,dtype,block,num_warps,launches,warmup,repeat,seed,process,scope,input,sample,ms\n";
    file << std::setprecision(17);
    for (std::size_t i = 0; i < timing.samples.size(); ++i)
        file << args.config_id << ',' << implementation.name << ",hip,"
             << args.rows << 'x' << args.columns << ",float32," << args.block
             << ",0," << implementation.launches << ',' << args.warmup << ','
             << args.repeat << ',' << args.seed << ',' << args.process << ",gpu-event-full-operator,shifted-dyadic,"
             << i << ',' << timing.samples[i] << '\n';
    if (!file) usage_error("failed writing sample file");
}

}  // namespace

int main(int argc, char** argv) {
    const Args args = parse_args(argc, argv);
    const std::size_t elements = args.rows * args.columns;
    const std::size_t bytes = elements * sizeof(float);
    const std::size_t row_bytes = args.rows * sizeof(float);

    int device = 0;
    hipDeviceProp_t properties{};
    HIP_CHECK(hipGetDevice(&device));
    HIP_CHECK(hipGetDeviceProperties(&properties, device));
    std::printf("ENV runtime=hip gpu=\"%s\" arch=%s wave_size=%d rows=%zu cols=%zu\n",
                properties.name, properties.gcnArchName, properties.warpSize, args.rows, args.columns);

    const std::vector<float> host_input = make_input(args);
    const std::vector<float> reference =
        cpu_reference(host_input, args.rows, args.columns);
    std::vector<float> host_output(elements);

    float* device_input = nullptr;
    float* device_output = nullptr;
    float* device_exponentials = nullptr;
    float* device_row_max = nullptr;
    float* device_row_sum = nullptr;
    HIP_CHECK(hipMalloc(&device_input, bytes));
    HIP_CHECK(hipMalloc(&device_output, bytes));
    HIP_CHECK(hipMalloc(&device_exponentials, bytes));
    HIP_CHECK(hipMalloc(&device_row_max, row_bytes));
    HIP_CHECK(hipMalloc(&device_row_sum, row_bytes));
    HIP_CHECK(hipMemcpy(device_input, host_input.data(), bytes,
                        hipMemcpyHostToDevice));

    bool all_correct = true;
    for (const Implementation& implementation :
         selected_implementations(args.version)) {
        HIP_CHECK(hipMemset(device_output, 0xFF, bytes));
        HIP_CHECK(hipMemset(device_exponentials, 0xFF, bytes));
        HIP_CHECK(hipMemset(device_row_max, 0xFF, row_bytes));
        HIP_CHECK(hipMemset(device_row_sum, 0xFF, row_bytes));
        launch(implementation, args, device_input, device_output,
               device_exponentials, device_row_max, device_row_sum);
        HIP_CHECK(hipDeviceSynchronize());
        HIP_CHECK(hipMemcpy(host_output.data(), device_output, bytes,
                            hipMemcpyDeviceToHost));
        const Validation precheck =
            validate(host_output, reference, args.rows, args.columns);
        if (!precheck.correct) {
            print_result(implementation, args, nullptr, precheck,
                         "FAIL", "NA");
            all_correct = false;
            continue;
        }

        const Timing timing = benchmark(
            implementation, args, device_input, device_output,
            device_exponentials, device_row_max, device_row_sum);
        HIP_CHECK(hipMemcpy(host_output.data(), device_output, bytes,
                            hipMemcpyDeviceToHost));
        const Validation postcheck =
            validate(host_output, reference, args.rows, args.columns);
        print_result(implementation, args, &timing, postcheck,
                     "OK", postcheck.correct ? "OK" : "FAIL");
        all_correct = all_correct && postcheck.correct;
        if (postcheck.correct) save_samples(implementation, args, timing);
    }

    HIP_CHECK(hipFree(device_row_sum));
    HIP_CHECK(hipFree(device_row_max));
    HIP_CHECK(hipFree(device_exponentials));
    HIP_CHECK(hipFree(device_output));
    HIP_CHECK(hipFree(device_input));
    return all_correct ? EXIT_SUCCESS : EXIT_FAILURE;
}

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

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

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

softmax_hip.hip 的 baseline 版本便于对照公式：前两个 kernel 各让一个线程处理一整行，第三个 kernel 再按元素写回。

kernel	分工	写出什么
row_max_serial	一个线程循环读取整行	row_max[row]
row_exp_sum_serial	同样逐列循环，取指数并求和	exponentials、row_sum[row]
normalize_rows	每个有效线程处理一个元素	exponentials[index] / row_sum[index / columns]

同一行通过全局数组交接

这里只跟踪一行，数组的原值在读取后保留。箭头是数据依赖，不表示耗时或硬件线路。

图 10.3　三次调用的静态交接图，只跟踪一行 [1000,1001,1002]。橙色框是全局中间数组：后续调用读取这些结果，输入和中间数组不会因读取消失。

图 10.3 中，第二次调用需要原输入和行最大值，第三次调用需要保存的指数和行和。先沿这些依赖确认中间数组的用途。进入优化之前，还要解决它的行内分工：一个线程独自扫描 1024 个数，同行没有用上多个线程的协作；同一轮中相邻线程还在读取不同行，地址相隔 columns 个元素。

先运行公式版本，检查逐元素误差和行和误差：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
test ! -e "$RUN/n4096x1024-hip-serial-b256-p1.csv" && \
"$RUN/build/softmax_hip" \
  --block 256 --version baseline --rows 4096 \
  --cols 1024 --warmup 10 --repeat 50 \
  --seed 20260920 --config-id hip-serial-b256 --samples "$RUN/n4096x1024-hip-serial-b256-p1.csv" \
  --process 1

真实输出：hip-serial-b256，shape=4096x1024，第 1 个独立进程

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 wave_size=32 rows=4096 cols=1024
RESULT operator=softmax implementation=hip-baseline-3kernel runtime=hip shape=4096x1024 dtype=float32 block=256 launches=3 lds_bytes=0 warmup=10 repeat=50 seed=20260920 timing=gpu-event scope=gpu-event-full-operator input=shifted-dyadic reference=fp64-softmax-to-fp32 timed=1 correct=OK precheck=OK postcheck=OK config_id=hip-serial-b256 num_warps=0 output_finite=OK min_ms=0.707899988 median_ms=0.765038997 mean_ms=0.765107559 max_abs_error=3.35276127e-08 max_row_sum_error=1.03569604e-06

先确认 correct/precheck/postcheck=OK，再核对 shape=4096x1024,launches=3。median_ms 是这个进程的完整三阶段时间，不能只选最短一轮作为成绩。

因此先保留三个阶段和这些中间数组，把每一行交给整个 block。这一步建立后面的融合基线，避免将行内并行化和阶段融合一次混在同一个比较里。

10.5.2 建立基线：每行一个 block，保留三个阶段

新版 cooperative 仍由三个 kernel 组成，每个 kernel 都让一个 block 处理一行。blockIdx.x 是行号，线程 t 负责列 t、t+blockDim.x、t+2×blockDim.x……。

先把 block 缩为 4 个线程、列数设为 10，就能检查分工：

block 内线程	负责的列
0	0、4、8
1	1、5、9
2	2、6
3	3、7

每列恰好处理一次，每一轮相邻线程读取相邻列。真实主配置使用 block=256、C=1024，所以每个线程先处理 4 列。

第一个阶段 row_max_block 先各自求局部最大值，再用第 9 章的 LDS 树合并：

const unsigned int lane = threadIdx.x;
const std::size_t base = static_cast<std::size_t>(blockIdx.x) * columns;
float local = kDeviceNegativeInfinity;
for (std::size_t column = lane; column < columns; column += blockDim.x)
    local = fmaxf(local, input[base + column]);
shared[lane] = local;
__syncthreads();
for (unsigned int stride = blockDim.x / 2; stride > 0; stride /= 2) {
    if (lane < stride) shared[lane] = fmaxf(shared[lane], shared[lane + stride]);
    __syncthreads();
}
if (lane == 0) row_max[blockIdx.x] = shared[0];

这里沿用源码中的变量名 lane，它表示 threadIdx.x，即 block 内线程编号。没有有效列的线程保持 max 的单位元，仍要到达 block 屏障。 第二阶段 row_exp_sum_block 使用同样的列分工，每个线程计算自己负责的指数、写入全局指数数组，并累加局部和：

const float maximum = row_max[blockIdx.x];
float local = 0.0F;
for (std::size_t column = lane; column < columns; column += blockDim.x) {
    const float value = expf(input[base + column] - maximum);
    exponentials[base + column] = value;
    local += value;
}

之后用 LDS 求和树得到 row_sum。第三阶段 normalize_block 仍按一行一个 block，把保存的指数除以这行的分母。三个阶段之间通过同一 stream 的调用顺序交接；每个阶段内部则用 block 屏障协调 LDS 读写。

建立这一步对照的目的，是先观察行内协作。 serial → cooperative 没有减少 kernel 数，也没有去掉中间数组，但改变了归约方式、线程与列的映射以及启动规模。它们的总体变化不能简化成某一项硬件成本。

在同一份源码中选择 cooperative；输入与计时参数不变：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
test ! -e "$RUN/n4096x1024-hip-cooperative-b256-p1.csv" && \
"$RUN/build/softmax_hip" \
  --block 256 --version cooperative --rows 4096 \
  --cols 1024 --warmup 10 --repeat 50 \
  --seed 20260920 --config-id hip-cooperative-b256 --samples "$RUN/n4096x1024-hip-cooperative-b256-p1.csv" \
  --process 1

真实输出：hip-cooperative-b256，shape=4096x1024，第 1 个独立进程

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 wave_size=32 rows=4096 cols=1024
RESULT operator=softmax implementation=hip-cooperative-3kernel runtime=hip shape=4096x1024 dtype=float32 block=256 launches=3 lds_bytes=1024 warmup=10 repeat=50 seed=20260920 timing=gpu-event scope=gpu-event-full-operator input=shifted-dyadic reference=fp64-softmax-to-fp32 timed=1 correct=OK precheck=OK postcheck=OK config_id=hip-cooperative-b256 num_warps=0 output_finite=OK min_ms=0.147676006 median_ms=0.149155997 mean_ms=0.160483680 max_abs_error=3.7252903e-09 max_row_sum_error=1.41068035e-07

同样先检查输出与行和，再比较完整时间。下面图表汇总三个进程，终端折叠区展示第 1 个进程。

两版均通过输出和行和检查。主输入的完整时间从 serial 的 773.261 μs 降至 cooperative 的 149.296 μs。启动代码中，串行版前两阶段各用 16 个 block，协作版三个阶段各用 4,096 个 block，每块 256 线程。分工可从代码核对；这项完整时间差没有单独测出某一种成本。

图 10.4　行内协作准备对照：两版保留相同中间数组和三个阶段，变化的是怎样将一行分配给线程。后续融合统一与 cooperative 比较。

本轮先采用 cooperative 作为 HIP 性能基线。它仍要保存整张指数数组，仍通过三个 kernel 完成。下一步围绕这些中间交接提出融合假设，而不是继续用串行版本作为所有收益的分母。

10.5.3 融合对照：保持行分工，减少中间交接

现在有了每行一个 block 的正确实现，三个阶段之间仍要传递 row_max、exponentials、row_sum。下一项假设是：保持这套行分工，在同一个 kernel 中完成 max、sum 和归一化，减少中间交接能否覆盖重读与重算的代价？

fused 保持 block=256，沿用相同的列步长和两棵 LDS 树。区别在于行最大值与分母在一次 kernel 内交接，不再通过全局行标量数组传递。指数不写入 exponentials，最后输出时重新计算。

LDS 复用必须等到所有读者完成

max 树结束后，行最大值位于 shared[0]。每个线程先把它读到自己的局部变量，再共同经过一道屏障：

const float maximum = shared[0];
// Every thread must read the maximum before LDS is reused for the sum.
__syncthreads();

前一轮屏障保证最大值已经写好；这里的屏障保证所有线程都已经读完。否则，一个 wavefront 先把 LDS 改作 sum 空间，另一个尚未读取最大值的 wavefront 就可能读到被覆盖的数据。

接下来每个线程再次扫描自己的列，累加 expf(input - maximum)；将局部和写入 LDS，用同样的树求出分母。最后再扫描一次自己的列，写出：

output[base + column] =
    expf(input[base + column] - maximum) / denominator;

HIP · 同一共享槽的两种用途读完最大值，才能覆盖共享槽

1/4

shared[0] 保存最大值，读者尚未取得副本沿用输入 [1000,1001,1002]，最大值树已经完成。这里只观察同一 block 的线程 0 和线程 32，它们属于不同 wave；其余 254 个线程省略。

读出副本，原来的 shared[0] 保持 1002两个代表线程先后取得 maximum=1002 的副本。画面结束时，本例整个 block 的读者都已读完；在此之前，任何线程都不能抢先将 shared[0] 改作 sum。

所有线程共同经过读者屏障全 block 到达这道 __syncthreads()，才能继续写下一轮共享数据。前面的树屏障保证最大值已写好，这道屏障保证所有线程已读完；它们保护的是不同的依赖。

线程 0 写局部 sum，maximum 的副本仍为 1002本例只有 3 列、block=256，因此线程 0 只负责第 0 列。它算出 exp(1000−1002)≈0.1353，再写到 shared[0]。共享槽的用途变了，各线程自己的 maximum 副本继续保留，供后续运算使用。

shared[0] 保存最大值，读者尚未取得副本沿用输入 [1000,1001,1002]，最大值树已经完成。这里只观察同一 block 的线程 0 和线程 32，它们属于不同 wave；其余 254 个线程省略。

步骤 1 / 4最大值就位

只跟踪 shared[0] 与两个代表读者；其余线程也须经过全 block 屏障。副本表示局部变量，动画节奏不是实际调度或耗时。

图 10.5　LDS 复用特写：shared[0] 的 1002 被读成各线程自己的 maximum 副本；全 block 经过读者屏障后，线程 0 才将局部指数和约 0.1353 写入这个共享槽。本例采用 3 列、256 线程，只显示两个代表读者。

在 图 10.5 的“允许覆盖”一步，比较共享槽和两个 maximum 副本：只有共享槽变成约 0.1353，已经取得的副本仍为 1002。线程 32 没有有效列，也须读最大值并参与 block 屏障；图外其余线程遵循相同规则。无有效列的线程给 max 贡献负无穷，给 sum 贡献 0，不能跳过同步。

把减少和增加的工作一起列出来

融合没有把输入变成只读一次。按当前源码，每个有效元素在 max、sum、最终输出三个阶段各被读取一次，指数在 sum 和输出时各计算一次：

源码中的工作	cooperative 三阶段基线	fused
每次完整调用的 kernel 数	3	1
输入 X 的扫描次数	2	3
每个元素的 expf 次数	1	2
指数数组	写一次、读一次	不读写
row_max、row_sum	经全局数组传递	在一次 kernel 内使用
行内分工与归约	一行一 block、列步长相同、LDS 树	保持同一方案，并正确同步 LDS 复用

主输入每张 FP32 矩阵是 16 MiB。只数输入、指数数组和输出这三类大数组，基线为 2×读X + 写exp + 读exp + 写Y，共 80 MiB；融合版为 3×读X + 写Y，共 64 MiB。去掉指数数组的 32 MiB 往返，同时新增一遍 16 MiB 输入读取，还多算了一次指数。

这是按源码估算的逻辑访问量，另有行标量和 LDS 访问；缓存与编译器会影响实际硬件流量。当前统一测试程序还为不同版本分配了相同的中间缓冲区，所以这里说明的是 fused 不再读写这些数组，不能直接当作整个进程峰值显存减少的证明。

运行 fused，完整输入和检查标准与 cooperative 相同：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
test ! -e "$RUN/n4096x1024-hip-fused-b256-p1.csv" && \
"$RUN/build/softmax_hip" \
  --block 256 --version fused --rows 4096 \
  --cols 1024 --warmup 10 --repeat 50 \
  --seed 20260920 --config-id hip-fused-b256 --samples "$RUN/n4096x1024-hip-fused-b256-p1.csv" \
  --process 1

真实输出：hip-fused-b256，shape=4096x1024，第 1 个独立进程

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 wave_size=32 rows=4096 cols=1024
RESULT operator=softmax implementation=hip-fused-block-lds runtime=hip shape=4096x1024 dtype=float32 block=256 launches=1 lds_bytes=1024 warmup=10 repeat=50 seed=20260920 timing=gpu-event scope=gpu-event-full-operator input=shifted-dyadic reference=fp64-softmax-to-fp32 timed=1 correct=OK precheck=OK postcheck=OK config_id=hip-fused-b256 num_warps=0 output_finite=OK min_ms=0.118597001 median_ms=0.119717002 mean_ms=0.121076540 max_abs_error=3.7252903e-09 max_row_sum_error=1.41068035e-07

检查两种误差与前后校验，再读 launches=1 和完整时间。是否采用由这组不带 profiler 的 benchmark 决定。

图 10.6　保持行分工后比较融合方案。图中时间包含得到整张输出的完整路径，不能只取其中一个归约阶段。

完整时间由 149.296 μs 降至 119.717 μs，约为基线时间的 80.2%，速度比约 1.25×。因此对主输入保留 fused。这支持的是减少交接、增加重读与重算之后的净收益，不能把约 29.6 μs 全部归为某次内存访问或某两次启动。

按需分析：融合是否按计划保持行分工、减少阶段交接？

已经知道 fused 更快，下面只核对执行结构。仍从 code/part2-kernels/ 运行，对 cooperative 与 fused 分别采集；每条命令新建自己的目录，已有目录会停止本条命令。随后用第 8 章的标准库脚本按 kernel 名分开读：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
mkdir "$RUN/profile-hip-cooperative-b256" && \
rocprofv3 --kernel-trace --output-format csv \
  --output-directory "$RUN/profile-hip-cooperative-b256" --output-file hip-cooperative-b256 -- \
  "$RUN/build/softmax_hip" \
    --block 256 --version cooperative --rows 4096 \
    --cols 1024 --warmup 5 --repeat 10 \
    --seed 20260920 --config-id hip-cooperative-b256 && \
python chapter8/inspect_trace.py \
  "$RUN/profile-hip-cooperative-b256/hip-cooperative-b256_kernel_trace.csv" \
  --kernel row_max_block --skip 6 --take 10 && \
python chapter8/inspect_trace.py \
  "$RUN/profile-hip-cooperative-b256/hip-cooperative-b256_kernel_trace.csv" \
  --kernel row_exp_sum_block --skip 6 --take 10 && \
python chapter8/inspect_trace.py \
  "$RUN/profile-hip-cooperative-b256/hip-cooperative-b256_kernel_trace.csv" \
  --kernel normalize_block --skip 6 --take 10

筛选后的真实 trace：hip-cooperative-b256，shape=4096x1024

kernel_name="(anonymous namespace)::row_max_block(float const*, float*, unsigned long)"
csv_rows=52 target_rows=16 skipped_rows=6 kept_rows=10
sort=Start_Timestamp_ascending
first_start_ns=885352907365890 first_end_ns=885352907410649 first_duration_ns=44759
kernel_min_us=44.519000 kernel_median_us=44.759000 kernel_max_us=270.714000
grid_work_items_xyz=1048576x1x1 workgroup_size_xyz=256x1x1 workgroups_xyz=4096x1x1
workgroups_total=4096
vgpr_count=8 sgpr_count=128 lds_bytes=0 scratch_bytes=0

kernel_name="(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)"
csv_rows=52 target_rows=16 skipped_rows=6 kept_rows=10
sort=Start_Timestamp_ascending
first_start_ns=885352907414689 first_end_ns=885352907475607 first_duration_ns=60918
kernel_min_us=56.679000 kernel_median_us=60.158500 kernel_max_us=60.918000
grid_work_items_xyz=1048576x1x1 workgroup_size_xyz=256x1x1 workgroups_xyz=4096x1x1
workgroups_total=4096
vgpr_count=16 sgpr_count=128 lds_bytes=0 scratch_bytes=0

kernel_name="(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)"
csv_rows=52 target_rows=16 skipped_rows=6 kept_rows=10
sort=Start_Timestamp_ascending
first_start_ns=885352907481287 first_end_ns=885352907541366 first_duration_ns=60079
kernel_min_us=36.439000 kernel_median_us=59.158500 kernel_max_us=61.118000
grid_work_items_xyz=1048576x1x1 workgroup_size_xyz=256x1x1 workgroups_xyz=4096x1x1
workgroups_total=4096
vgpr_count=16 sgpr_count=128 lds_bytes=0 scratch_bytes=0

原始 kernel trace CSV：hip-cooperative-b256，shape=4096x1024

"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,1200097,1,7,"__amd_rocclr_fillBufferUnAligned",1,885352876454435,885352876478714,0,0,8,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,2,7,"__amd_rocclr_fillBufferUnAligned",2,885352876642399,885352876657759,0,0,8,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,3,7,"__amd_rocclr_fillBufferUnAligned",3,885352876669839,885352876672199,0,0,8,0,128,256,1,1,1024,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,4,7,"__amd_rocclr_fillBufferUnAligned",4,885352876677238,885352876680158,0,0,8,0,128,256,1,1,1024,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,5,19,"(anonymous namespace)::row_max_block(float const*, float*, unsigned long)",5,885352888391880,885352888442878,0,0,8,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,6,20,"(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)",6,885352888448118,885352888514837,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,7,23,"(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)",7,885352888523237,885352888583235,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,8,19,"(anonymous namespace)::row_max_block(float const*, float*, unsigned long)",8,885352906389885,885352906440004,0,0,8,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,9,20,"(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)",9,885352906445124,885352906511681,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,10,23,"(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)",10,885352906519761,885352906581639,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,11,19,"(anonymous namespace)::row_max_block(float const*, float*, unsigned long)",11,885352906589719,885352906639558,0,0,8,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,18,20,"(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)",18,885352907016149,885352907076947,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,13,23,"(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)",13,885352906717316,885352906778834,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,14,19,"(anonymous namespace)::row_max_block(float const*, float*, unsigned long)",14,885352906784994,885352906829713,0,0,8,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,15,20,"(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)",15,885352906834033,885352906894912,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,19,23,"(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)",19,885352907083027,885352907142185,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,16,23,"(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)",16,885352906900951,885352906960790,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,12,20,"(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)",12,885352906644358,885352906710596,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,17,19,"(anonymous namespace)::row_max_block(float const*, float*, unsigned long)",17,885352906966910,885352907011669,0,0,8,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,20,19,"(anonymous namespace)::row_max_block(float const*, float*, unsigned long)",20,885352907148305,885352907193144,0,0,8,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,21,20,"(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)",21,885352907197544,885352907258703,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,22,23,"(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)",22,885352907264822,885352907323461,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,23,19,"(anonymous namespace)::row_max_block(float const*, float*, unsigned long)",23,885352907365890,885352907410649,0,0,8,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,24,20,"(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)",24,885352907414689,885352907475607,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,25,23,"(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)",25,885352907481287,885352907541366,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,26,19,"(anonymous namespace)::row_max_block(float const*, float*, unsigned long)",26,885352907577925,885352907622604,0,0,8,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,27,20,"(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)",27,885352907626724,885352907686922,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,28,23,"(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)",28,885352907692722,885352907751801,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,29,19,"(anonymous namespace)::row_max_block(float const*, float*, unsigned long)",29,885352907782720,885352907827719,0,0,8,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,30,20,"(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)",30,885352907831759,885352907892197,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,31,23,"(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)",31,885352907897957,885352907959075,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,32,19,"(anonymous namespace)::row_max_block(float const*, float*, unsigned long)",32,885352907990195,885352908034954,0,0,8,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,33,20,"(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)",33,885352908039033,885352908099832,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,34,23,"(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)",34,885352908105552,885352908165990,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,35,19,"(anonymous namespace)::row_max_block(float const*, float*, unsigned long)",35,885352908197110,885352908241868,0,0,8,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,36,20,"(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)",36,885352908245988,885352908306307,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,37,23,"(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)",37,885352908312027,885352908371265,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,38,19,"(anonymous namespace)::row_max_block(float const*, float*, unsigned long)",38,885352908402345,885352908447263,0,0,8,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,39,20,"(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)",39,885352908451383,885352908511502,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,40,23,"(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)",40,885352908517302,885352908577420,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,41,19,"(anonymous namespace)::row_max_block(float const*, float*, unsigned long)",41,885352908608819,885352908879533,0,0,8,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,42,20,"(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)",42,885352908883693,885352908940571,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,43,23,"(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)",43,885352908946291,885352908984890,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,44,19,"(anonymous namespace)::row_max_block(float const*, float*, unsigned long)",44,885352909016009,885352909060568,0,0,8,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,45,20,"(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)",45,885352920373090,885352920431369,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,46,23,"(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)",46,885352920437809,885352920475288,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,47,19,"(anonymous namespace)::row_max_block(float const*, float*, unsigned long)",47,885352920514967,885352920561046,0,0,8,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,48,20,"(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)",48,885352920565286,885352920622484,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,49,23,"(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)",49,885352920628324,885352920664763,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,50,19,"(anonymous namespace)::row_max_block(float const*, float*, unsigned long)",50,885352920700802,885352920745321,0,0,8,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,51,20,"(anonymous namespace)::row_exp_sum_block(float const*, float const*, float*, float*, unsigned long)",51,885352920749441,885352920806120,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200097,52,23,"(anonymous namespace)::normalize_block(float const*, float const*, float*, unsigned long)",52,885352920811960,885352920848599,0,0,16,0,128,256,1,1,1048576,1,1

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
mkdir "$RUN/profile-hip-fused-b256" && \
rocprofv3 --kernel-trace --output-format csv \
  --output-directory "$RUN/profile-hip-fused-b256" --output-file hip-fused-b256 -- \
  "$RUN/build/softmax_hip" \
    --block 256 --version fused --rows 4096 \
    --cols 1024 --warmup 5 --repeat 10 \
    --seed 20260920 --config-id hip-fused-b256 && \
python chapter8/inspect_trace.py \
  "$RUN/profile-hip-fused-b256/hip-fused-b256_kernel_trace.csv" \
  --kernel softmax_fused_lds --skip 6 --take 10

筛选后的真实 trace：hip-fused-b256，shape=4096x1024

kernel_name="(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)"
csv_rows=20 target_rows=16 skipped_rows=6 kept_rows=10
sort=Start_Timestamp_ascending
first_start_ns=885353157818181 first_end_ns=885353157928658 first_duration_ns=110477
kernel_min_us=110.117000 kernel_median_us=110.577500 kernel_max_us=118.757000
grid_work_items_xyz=1048576x1x1 workgroup_size_xyz=256x1x1 workgroups_xyz=4096x1x1
workgroups_total=4096
vgpr_count=16 sgpr_count=128 lds_bytes=0 scratch_bytes=0

原始 kernel trace CSV：hip-fused-b256，shape=4096x1024

"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,1200117,1,5,"__amd_rocclr_fillBufferUnAligned",1,885353127709195,885353127724394,0,0,8,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,2,5,"__amd_rocclr_fillBufferUnAligned",2,885353127893189,885353127903189,0,0,8,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,3,5,"__amd_rocclr_fillBufferUnAligned",3,885353127911869,885353127914149,0,0,8,0,128,256,1,1,1024,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,4,5,"__amd_rocclr_fillBufferUnAligned",4,885353127918629,885353127921229,0,0,8,0,128,256,1,1,1024,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,5,22,"(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)",5,885353139218459,885353139332616,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,6,22,"(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)",6,885353157179919,885353157296556,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,7,22,"(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)",7,885353157300996,885353157414192,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,8,22,"(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)",8,885353157420191,885353157539668,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,9,22,"(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)",9,885353157545508,885353157656425,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,10,22,"(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)",10,885353157662265,885353157772942,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,11,22,"(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)",11,885353157818181,885353157928658,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,12,22,"(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)",12,885353157963737,885353158082494,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,13,22,"(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)",13,885353158114734,885353158225331,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,14,22,"(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)",14,885353158257450,885353158367927,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,15,22,"(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)",15,885353158399646,885353158510563,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,16,22,"(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)",16,885353158541642,885353158652400,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,17,22,"(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)",17,885353158683279,885353158793916,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,18,22,"(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)",18,885353158824795,885353158935152,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,19,22,"(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)",19,885353158965791,885353159076349,0,0,16,0,128,256,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",1,0,1200117,20,22,"(anonymous namespace)::softmax_fused_lds(float const*, float*, unsigned long, unsigned long)",20,885353159107028,885353159217145,0,0,16,0,128,256,1,1,1048576,1,1

先读 target_rows=16，排除预检和预热的前 6 次，留下 10 次。再看 kernel 名称：cooperative 的三个阶段与 fused 的一个阶段分别存在；不能把整个文件的行数当作一次调用的 kernel 数。四个阶段的 workgroups_total=4096、每组 256 线程，符合每行一个 block 的对照设计。

最后才看 kernel_median_us：fused 的设备区间约 110.6 μs。它不与另一轮 event 相减估算启动成本，也不把 cooperative 三个独立中位数之和作为正式成绩。首阶段有约 270.7 μs 的单次值，报告保留完整范围。

这些记录支持“融合遵循了预期的行分工与阶段数量”，仍不能把总收益分给某次读取或某次启动。HIP 的 lds_bytes=0 未包含这次动态 LDS 申请，源码实际请求 256×4=1024 B，不能读成没有 LDS。

10.5.4 检查行长与行数的影响

一个主形状不足以决定今后的配置。先保持行数不变，检查列数 1023、1024、1025：对于 block=256，1024 列时每线程处理 4 项；1025 列时线程 0 多处理一项，其他线程的分工基本不变。1023 列则有一个线程少处理一项，尾部仍要正确参与两次归约。

先用一个小尾部验证 fused 的分工。此处只检查答案，warmup=0,repeat=1 的时间不进入性能图：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
"$RUN/build/softmax_hip" \
  --block 256 --version fused --rows 3 \
  --cols 1025 --warmup 0 --repeat 1 \
  --seed 20260920 --config-id hip-fused-b256

正确性输出：HIP fused，3×1025，输出与行和

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 wave_size=32 rows=3 cols=1025
RESULT operator=softmax implementation=hip-fused-block-lds runtime=hip shape=3x1025 dtype=float32 block=256 launches=1 lds_bytes=1024 warmup=0 repeat=1 seed=20260920 timing=gpu-event scope=gpu-event-full-operator input=shifted-dyadic reference=fp64-softmax-to-fp32 timed=1 correct=OK precheck=OK postcheck=OK config_id=hip-fused-b256 num_warps=0 output_finite=OK min_ms=0.016480001 median_ms=0.016480001 mean_ms=0.016480001 max_abs_error=3.7252903e-09 max_row_sum_error=2.65218175e-08

先读前后校验、有限值与两种误差；尾部正确后，才比较较大输入的性能。

再保持 1024 列，把行数从 4096 减少到 128。每行一个 block 意味着 grid 也从 4096 个 block 变为 128 个；可供 GPU 调度的组数和总工作量都减少了。阶段启动、中间数组与线程协作的相对成本可能随形状变化，因此要在新形状内重新比较 cooperative 和 fused。

形状	cooperative（μs）	fused（μs）
4096x1023	150.40（147.12–151.52）	120.38（120.06–120.48）
4096x1024	149.30（149.16–149.96）	119.72（119.72–119.74）
4096x1025	156.98（153.76–157.86）	130.52（130.26–130.62）
128x1024	24.68（24.68–24.80）	15.88（15.84–15.90）

括号表示三进程 median 范围；每行在相同形状内比较。

三个补充形状里，fused 也都低于同形状的 cooperative。主形状的独立确认中，cooperative 为 146.857 μs（145.938–149.297），fused 为 119.858 μs（117.098–119.858），这项收益再次出现。确认组另行统计，不并入原来的图。这里没有扫描其他 block 大小，因此结论是保留本轮融合方案，并非找到了每种形状的最佳实现。

10.5.5 Triton 基线：一行一个 program

完整源码：code/part2-kernels/chapter10/softmax_triton.py

"""Chapter 10: stable row-wise Softmax, one Triton program per row.

Performance numbers from this script belong in the tutorial only after a run
on the documented ROCm machine. The host code checks a CPU FP64 Softmax reference rounded to FP32
before and after timing and prints machine-readable RESULT lines.
"""

from __future__ import annotations

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

import torch
import triton
import triton.language as tl


ABSOLUTE_TOLERANCE = 2.0e-5
ROW_SUM_TOLERANCE = 2.0e-5
MAX_BLOCK_SIZE = 65536


@triton.jit
def softmax_row_kernel(
    input_ptr,
    output_ptr,
    columns,
    input_row_stride,
    output_row_stride,
    BLOCK_SIZE: tl.constexpr,
):
    """Compute one complete row with a numerically stable Softmax."""
    row = tl.program_id(axis=0)
    offsets = tl.arange(0, BLOCK_SIZE)
    valid = offsets < columns

    logits = tl.load(
        input_ptr + row * input_row_stride + offsets,
        mask=valid,
        other=-float("inf"),
    )
    maximum = tl.max(logits, axis=0)
    numerator = tl.exp(logits - maximum)
    denominator = tl.sum(numerator, axis=0)
    probabilities = numerator / denominator
    tl.store(
        output_ptr + row * output_row_stride + offsets,
        probabilities,
        mask=valid,
    )


@dataclass(frozen=True)
class Implementation:
    name: str
    block_size: int
    num_warps: int


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


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


def next_power_of_two(value: int) -> int:
    return 1 << (value - 1).bit_length()


def parse_args() -> argparse.Namespace:
    parser = argparse.ArgumentParser()
    parser.add_argument("--version", choices=("all", "t0", "t1"), default="all")
    parser.add_argument("--rows", type=int, default=4096)
    parser.add_argument("--cols", type=int, default=1024)
    parser.add_argument("--t0-warps", type=int, default=4)
    parser.add_argument(
        "--t1-block",
        type=int,
        default=0,
        help="0 chooses twice t0's block when possible",
    )
    parser.add_argument("--t1-warps", type=int, default=8)
    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 args.rows <= 0 or args.cols <= 0:
        parser.error("--rows and --cols must be positive")
    if args.cols > MAX_BLOCK_SIZE:
        parser.error(f"--cols must be at most {MAX_BLOCK_SIZE}")
    if args.t1_block < 0:
        parser.error("--t1-block must be zero or positive")
    if args.t1_block and args.t1_block & (args.t1_block - 1):
        parser.error("--t1-block must be a power of two")
    if args.t1_block and args.t1_block < args.cols:
        parser.error("--t1-block must cover --cols")
    if args.t1_block > MAX_BLOCK_SIZE:
        parser.error(f"--t1-block must be at most {MAX_BLOCK_SIZE}")
    if (
        args.t0_warps not in (1, 2, 4, 8)
        or args.t1_warps not in (1, 2, 4, 8)
    ):
        parser.error("warp counts must be one of 1, 2, 4, or 8")
    if args.warmup < 0 or args.repeat <= 0:
        parser.error("--warmup must be non-negative and --repeat positive")
    if args.samples and (args.version == "all" or not args.config_id):
        parser.error("--samples requires one version and a --config-id")
    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.process <= 0:
        parser.error("--process must be positive")
    return args


def selected_implementations(args: argparse.Namespace) -> list[Implementation]:
    t0_block = next_power_of_two(args.cols)
    automatic_t1_block = min(t0_block * 2, MAX_BLOCK_SIZE)
    t1_block = args.t1_block or automatic_t1_block
    implementations = {
        "t0": Implementation("triton-t0-compact", t0_block, args.t0_warps),
        "t1": Implementation("triton-t1-wide", t1_block, args.t1_warps),
    }
    if args.version == "all":
        return [implementations["t0"], implementations["t1"]]
    return [implementations[args.version]]


def make_host_input(rows: int, columns: int, seed: int) -> torch.Tensor:
    """Match the HIP program's deterministic stability-stress input."""
    row_ids = torch.arange(rows, dtype=torch.int64).view(rows, 1)
    column_ids = torch.arange(columns, dtype=torch.int64).view(1, columns)
    mixed = column_ids * 17 + row_ids * 131 + seed
    centered = torch.remainder(mixed, 4096) - 2048
    shifts = torch.where(
        torch.remainder(row_ids, 3) == 0,
        torch.tensor(1000.0, dtype=torch.float32),
        torch.where(
            torch.remainder(row_ids, 3) == 1,
            torch.tensor(-1000.0, dtype=torch.float32),
            torch.tensor(0.0, dtype=torch.float32),
        ),
    )
    return centered.to(torch.float32) / 128.0 + shifts


def launch(
    implementation: Implementation,
    input_tensor: torch.Tensor,
    output_tensor: torch.Tensor,
) -> None:
    rows, columns = input_tensor.shape
    softmax_row_kernel[(rows,)](
        input_tensor,
        output_tensor,
        columns,
        input_tensor.stride(0),
        output_tensor.stride(0),
        BLOCK_SIZE=implementation.block_size,
        num_warps=implementation.num_warps,
    )


def validate(output: torch.Tensor, reference: torch.Tensor) -> Validation:
    actual = output.detach().to(device="cpu")
    difference = torch.abs(actual - reference)
    max_absolute_error = float(difference.amax().item())
    max_row_sum_error = float(
        torch.abs(actual.double().sum(dim=1) - 1.0).amax().item()
    )
    finite = bool(torch.isfinite(actual).all().item())
    return Validation(
        correct=(
            finite
            and max_absolute_error <= ABSOLUTE_TOLERANCE
            and max_row_sum_error <= ROW_SUM_TOLERANCE
        ),
        finite=finite,
        max_absolute_error=max_absolute_error,
        max_row_sum_error=max_row_sum_error,
    )


def benchmark(
    implementation: Implementation,
    input_tensor: torch.Tensor,
    output_tensor: torch.Tensor,
    *,
    warmup: int,
    repeat: int,
) -> Timing:
    for _ in range(warmup):
        launch(implementation, input_tensor, output_tensor)
    torch.cuda.synchronize()

    times_ms: list[float] = []
    for _ in range(repeat):
        start = torch.cuda.Event(enable_timing=True)
        end = torch.cuda.Event(enable_timing=True)
        start.record()
        launch(implementation, input_tensor, output_tensor)
        end.record()
        end.synchronize()
        times_ms.append(float(start.elapsed_time(end)))

    return Timing(
        minimum_ms=min(times_ms),
        median_ms=statistics.median(times_ms),
        mean_ms=statistics.fmean(times_ms),
        samples=tuple(times_ms),
    )


def print_result(
    implementation: Implementation,
    args: argparse.Namespace,
    validation: Validation,
    *,
    timing: Timing | None,
    precheck: str,
    postcheck: str,
) -> None:
    status = "OK" if validation.correct else "FAIL"
    timing_fields = (
        "min_ms=NA median_ms=NA mean_ms=NA"
        if timing is None
        else (
            f"min_ms={timing.minimum_ms:.6f} "
            f"median_ms={timing.median_ms:.6f} "
            f"mean_ms={timing.mean_ms:.6f}"
        )
    )
    print(
        "RESULT"
        " operator=softmax"
        f" implementation={implementation.name}"
        " runtime=triton"
        f" shape={args.rows}x{args.cols}"
        " dtype=float32"
        f" block={implementation.block_size}"
        f" num_warps={implementation.num_warps}"
        " programs_per_row=1"
        " launches=1"
        f" warmup={args.warmup}"
        f" repeat={args.repeat}"
        f" seed={args.seed}"
        " timing=gpu-event scope=gpu-event-full-operator input=shifted-dyadic reference=fp64-softmax-to-fp32"
        f" config_id={args.config_id}"
        f" output_finite={'OK' if validation.finite else 'FAIL'}"
        f" timed={int(timing is not None)}"
        f" correct={status}"
        f" precheck={precheck}"
        f" postcheck={postcheck}"
        f" {timing_fields}"
        f" max_abs_error={validation.max_absolute_error:.9g}"
        f" max_row_sum_error={validation.max_row_sum_error:.9g}"
    )


def main() -> int:
    args = parse_args()
    if not torch.cuda.is_available():
        raise RuntimeError("PyTorch cannot see a ROCm GPU")

    print(
        "ENV"
        " runtime=triton"
        f' gpu="{torch.cuda.get_device_name(0)}"'
        f" torch={torch.__version__}"
        f" torch_hip={torch.version.hip}"
        f" triton={triton.__version__}"
    )

    host_input = make_host_input(args.rows, args.cols, args.seed)
    reference = torch.softmax(host_input.double(), dim=1).float()
    device_input = host_input.to(device="cuda")
    output = torch.empty_like(device_input)

    all_correct = True
    for implementation in selected_implementations(args):
        output.fill_(float("nan"))
        launch(implementation, device_input, output)
        torch.cuda.synchronize()
        precheck = validate(output, reference)
        if not precheck.correct:
            print_result(
                implementation,
                args,
                precheck,
                timing=None,
                precheck="FAIL",
                postcheck="NA",
            )
            all_correct = False
            continue

        timing = benchmark(
            implementation,
            device_input,
            output,
            warmup=args.warmup,
            repeat=args.repeat,
        )
        postcheck = validate(output, reference)
        print_result(
            implementation,
            args,
            postcheck,
            timing=timing,
            precheck="OK",
            postcheck="OK" if postcheck.correct else "FAIL",
        )
        all_correct = all_correct and postcheck.correct
        if postcheck.correct and args.samples:
            with args.samples.open("w", newline="") as file:
                writer = csv.writer(file)
                writer.writerow(("config_id", "implementation", "runtime", "shape", "dtype",
                                 "block", "num_warps", "launches", "warmup", "repeat", "seed",
                                 "process", "scope", "input", "sample", "ms"))
                for sample, ms in enumerate(timing.samples):
                    writer.writerow((args.config_id, implementation.name, "triton",
                                     f"{args.rows}x{args.cols}", "float32", implementation.block_size,
                                     implementation.num_warps, 1, args.warmup, args.repeat, args.seed,
                                     args.process, "gpu-event-full-operator", "shifted-dyadic", sample, ms))

    return 0 if all_correct else 1


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

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

Triton 路线直接用一个逻辑向量表示一行。softmax_row_kernel 中的步骤与手算一一对应：

row = tl.program_id(axis=0)
offsets = tl.arange(0, BLOCK_SIZE)
valid = offsets < columns

logits = tl.load(
    input_ptr + row * input_row_stride + offsets,
    mask=valid,
    other=-float("inf"),
)
maximum = tl.max(logits, axis=0)
numerator = tl.exp(logits - maximum)
denominator = tl.sum(numerator, axis=0)
probabilities = numerator / denominator
tl.store(
    output_ptr + row * output_row_stride + offsets,
    probabilities,
    mask=valid,
)

Host 的 grid 为 (rows,)，每行启动一个 program。tl.max 和 tl.sum 都沿这一行的逻辑向量归约，其他行不会混进来。program 内保留了 numerator 这一逻辑值，没有显式的全局指数数组。

先让补齐位置符合数学

若有 3 列，先选能覆盖一行的最小 2 的幂宽度 B=4。用一行负数 [-1002,-1001,-1000]，更容易检查第四个逻辑位置应该怎样处理：

Triton · 一行一个 program补齐逻辑向量，保持原来三列的答案

1/5

3 个真实列，用 4 个逻辑位置表示输入为 −1002、−1001、−1000。program 的逻辑向量宽度 B=4，第四个位置没有对应的输入元素。这些格子是逻辑位置，不是四个物理线程。

有效位置读取输入，无效位置填负无穷前三个位置取得真实输入的副本；第四个位置的 load mask 为假，不读取行外地址，使用 other=−∞。原输入仍保留。

负无穷不改变这一行的最大值四个逻辑值共同做 max，结果仍为 −1000。若无效位置填 0，它反而会胜过所有真实的负输入；所以 load 的补值必须符合归约语义。

无效位置的指数为 0减去 −1000 后，逻辑值为 −2、−1、0、−∞；取指数得到约 0.1353、0.3679、1、0。第四项给 sum 贡献 0，分母仍为约 1.5032。

只有前三个结果写入输出前三个指数分别除以同一个分母，得到约 0.0900、0.2447、0.6652。store mask 仍只允许前三个位置写入；第四个逻辑位置不会创建一个额外输出。

3 个真实列，用 4 个逻辑位置表示输入为 −1002、−1001、−1000。program 的逻辑向量宽度 B=4，第四个位置没有对应的输入元素。这些格子是逻辑位置，不是四个物理线程。

步骤 1 / 5输入与位置

负数行手算示例。B 是逻辑宽度，具体线程映射由编译器安排；动画时间不代表 GPU 耗时。

图 10.7　真实输入只有三项。无效的第四个逻辑位置加载为 −∞，取指数后贡献 0；最后只写回三个真实输出。

在 图 10.7 中，停在“最大值”检查结果仍是 −1000，再看第四个逻辑位置怎样从 −∞ 变成指数 0。load mask 负责安全读取并提供正确补值，store mask 负责不写出行外。若加载时用 0 补齐，它会胜过所有真实负输入，把最大值改成 0；真实三项的指数可能全部下溢，结果就被破坏了。

基线选择 BLOCK_SIZE=B、num_warps=4，其中 B 是覆盖 C 的最小 2 的幂。对于主输入 C=1024，基线就是 (1024,4)。一个 program 有 1024 个逻辑位置，并不意味着它有 1024 个显式线程；编译器负责具体映射。

先运行紧凑宽度、4 warps 的基线：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
test ! -e "$RUN/n4096x1024-triton-b4-p1.csv" && \
python chapter10/softmax_triton.py \
  --t0-warps 4 --t1-warps 4 --t1-block 1024 \
  --version t0 --rows 4096 --cols 1024 \
  --warmup 10 --repeat 50 --seed 20260920 \
  --config-id triton-b4 --samples "$RUN/n4096x1024-triton-b4-p1.csv" --process 1

真实输出：triton-b4，shape=4096x1024，第 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=softmax implementation=triton-t0-compact runtime=triton shape=4096x1024 dtype=float32 block=1024 num_warps=4 programs_per_row=1 launches=1 warmup=10 repeat=50 seed=20260920 timing=gpu-event scope=gpu-event-full-operator input=shifted-dyadic reference=fp64-softmax-to-fp32 config_id=triton-b4 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.034399 median_ms=0.037199 mean_ms=0.048034 max_abs_error=5.58793545e-09 max_row_sum_error=1.20831897e-07

检查输出与行和，然后核对 block=1024、num_warps=4 与单次完整调用。基线三个进程汇总为 37.599 μs。这里逻辑宽度与硬件线程数是不同的量。

10.5.6 第一轮：固定逻辑宽度，比较 warps

基线已经用一个 program 完成一整行。接下来调整的不是 kernel 数，而是这个 program 的执行配置。先提出一个问题：同样长的一行，使用更多 warps 参与，能否缩短完整执行时间？

先保持 BLOCK_SIZE=B、输入、行数、kernel 代码不变，只比较 num_warps=4 与 8。两版仍是一行一个 program，读取相同的有效列，完成同样的 max、exp、sum 和归一化。

为检查这个选择是否依赖逻辑宽度，我们也准备 BLOCK_SIZE=2B 下的 4/8 对照。B 已经能覆盖整行，2B 只增加补齐位置，不增加有效列。因此图里可以分成两组，每组固定一个宽度，单独回答 warps 的变化。

更多 warps 会改变 program 内逻辑工作如何分配与协作，也可能改变资源使用。它不是“免费增加并行度”，所以先提出有限候选，再检查结果，不提前把 8 warps 当作优化版。

接下来只运行另外三项，不重复已有基线。四组数据会用于两个方向的比较：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
test ! -e "$RUN/n4096x1024-triton-b8-p1.csv" && \
python chapter10/softmax_triton.py \
  --t0-warps 8 --t1-warps 8 --t1-block 1024 \
  --version t0 --rows 4096 --cols 1024 \
  --warmup 10 --repeat 50 --seed 20260920 \
  --config-id triton-b8 --samples "$RUN/n4096x1024-triton-b8-p1.csv" --process 1

真实输出：triton-b8，shape=4096x1024，第 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=softmax implementation=triton-t0-compact runtime=triton shape=4096x1024 dtype=float32 block=1024 num_warps=8 programs_per_row=1 launches=1 warmup=10 repeat=50 seed=20260920 timing=gpu-event scope=gpu-event-full-operator input=shifted-dyadic reference=fp64-softmax-to-fp32 config_id=triton-b8 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.041079 median_ms=0.041439 mean_ms=0.050474 max_abs_error=3.7252903e-09 max_row_sum_error=1.2747008e-07

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
test ! -e "$RUN/n4096x1024-triton-2b4-p1.csv" && \
python chapter10/softmax_triton.py \
  --t0-warps 4 --t1-warps 4 --t1-block 2048 \
  --version t1 --rows 4096 --cols 1024 \
  --warmup 10 --repeat 50 --seed 20260920 \
  --config-id triton-2b4 --samples "$RUN/n4096x1024-triton-2b4-p1.csv" --process 1

真实输出：triton-2b4，shape=4096x1024，第 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=softmax implementation=triton-t1-wide runtime=triton shape=4096x1024 dtype=float32 block=2048 num_warps=4 programs_per_row=1 launches=1 warmup=10 repeat=50 seed=20260920 timing=gpu-event scope=gpu-event-full-operator input=shifted-dyadic reference=fp64-softmax-to-fp32 config_id=triton-2b4 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.051638 median_ms=0.051999 mean_ms=0.062111 max_abs_error=5.58793545e-09 max_row_sum_error=1.20831897e-07

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
test ! -e "$RUN/n4096x1024-triton-2b8-p1.csv" && \
python chapter10/softmax_triton.py \
  --t0-warps 8 --t1-warps 8 --t1-block 2048 \
  --version t1 --rows 4096 --cols 1024 \
  --warmup 10 --repeat 50 --seed 20260920 \
  --config-id triton-2b8 --samples "$RUN/n4096x1024-triton-2b8-p1.csv" --process 1

真实输出：triton-2b8，shape=4096x1024，第 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=softmax implementation=triton-t1-wide runtime=triton shape=4096x1024 dtype=float32 block=2048 num_warps=8 programs_per_row=1 launches=1 warmup=10 repeat=50 seed=20260920 timing=gpu-event scope=gpu-event-full-operator input=shifted-dyadic reference=fp64-softmax-to-fp32 config_id=triton-2b8 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.055998 median_ms=0.056279 mean_ms=0.061336 max_abs_error=3.7252903e-09 max_row_sum_error=1.2747008e-07

四项都先检查答案和行和，再核对逻辑宽度与 warps，最后比较 median_ms。同名 kernel 依靠不同 config-id 和样本路径区分；两张图复用这一组扫描，不是再采一轮。

图 10.8　每组固定一个逻辑宽度，单独比较 warps。原始 (B,4) 始终标出，先看紧凑宽度下的效果，再检查换一个宽度是否得到相同判断。

在紧凑宽度 1024 下，4→8 warps 使时间从 37.599 μs 增至 41.439 μs；在宽度 2048 下，也从 51.879 μs 增至 56.279 μs。两组都没有获益，因此主输入继续保留 4 warps。更多参与线程并没有自动缩短整行计算。

10.5.7 第二轮：固定 warps，比较逻辑宽度

现在从另一个方向读取这四项实验：固定 warps，判断把逻辑向量从 B 扩为 2B 是否值得。每次直接对照都只改一个参数：

配置	第一轮用途	第二轮用途
(B,4)	原始基线	固定 4 warps 的宽度基线
(B,8)	固定 B，增加 warps	固定 8 warps 的宽度基线
(2B,4)	—	保持 4 warps，只扩大逻辑宽度
(2B,8)	—	保持 8 warps，只扩大逻辑宽度

对于同一行，B 已足够覆盖全部有效列。扩大为 2B 没有新增有效输入，新增位置都被 mask。需要验证的是更宽逻辑向量带来的编译分配变化是否有利，还是只增加了无效位置与归约的成本；不能把这个参数解释为处理了更多有用数据。

源码层面，新增位置仍参与向量表达式，但其 load 为负无穷、指数为 0，store 被屏蔽。编译器如何映射或消除这些操作要看实际产物，因此不要简单将逻辑宽度翻倍换算成指令数或执行时间翻倍。

图 10.9　将两个参数拆开后，可以分别观察更改 warps 与扩大逻辑宽度的效果，也能检查宽度的效果是否依赖 warps 配置。

固定 4 warps 时，加宽使时间从 37.599 μs 增至 51.879 μs；固定 8 warps 时，则从 41.439 μs 增至 56.279 μs。所以主输入继续保留紧凑宽度。这个退步足以决定不采用加宽。若还要判断编译器增加了哪些资源需求，可再选固定 warps 的代表配置诊断；没有这项证据时，不把时间差命名为寄存器成本。

10.5.8 检查逻辑宽度边界与行数

Triton 的紧凑宽度随列数跳变：

实际列数 C	最小逻辑宽度 B	紧凑版无效位置	加宽版 2B	加宽版无效位置
1023	1024	1	2048	1025
1024	1024	0	2048	1024
1025	2048	1023	4096	3071

先确认新逻辑宽度没有污染 Softmax 的分母。下面使用同一份程序检查 3×1025，仅读正确性：

test -d "${RUN:?请先准备本章环境与结果目录}/build" && \
python chapter10/softmax_triton.py \
  --t0-warps 4 --t1-warps 4 --t1-block 2048 \
  --version t0 --rows 3 --cols 1025 \
  --warmup 0 --repeat 1 --seed 20260920 \
  --config-id triton-b4

正确性输出：Triton 紧凑宽度，3×1025

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=softmax implementation=triton-t0-compact runtime=triton shape=3x1025 dtype=float32 block=2048 num_warps=4 programs_per_row=1 launches=1 warmup=0 repeat=1 seed=20260920 timing=gpu-event scope=gpu-event-full-operator input=shifted-dyadic reference=fp64-softmax-to-fp32 config_id=triton-b4 output_finite=OK timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.020560 median_ms=0.020560 mean_ms=0.020560 max_abs_error=3.7252903e-09 max_row_sum_error=5.53661534e-08

这里 block=2048，但输出仍只有 1025 列；先看 correct 与行和误差，再看性能形状。

从 1024 列增加到 1025 列，输入只多一项，紧凑逻辑宽度却从 1024 变为 2048。因此即便时间出现明显变化，也不能只解释为“多算了一个数”。先核对实际 BLOCK_SIZE、warps 和 mask，再看编译与计时。

行数的影响也要单独检查：当前每行一个 program，所以行数决定了 program 总数；加宽一行不会增加 program 数。少量长行与大量同样长的行，可能需要不同的执行配置。

图 10.10　列数决定逻辑宽度，行数决定 program 数。图中先比较紧凑宽度的 4/8 warps；各面板刻度范围不同，不能跨面板比较柱长。全部四项配置均保存在发布记录中。

这次出现了配置顺序的变化：4096×1025 的紧凑宽度为 2048，4 warps 为 73.298 μs，8 warps 为 62.858 μs；此时 8 warps 更快。它与 4096×1024 的选择不同。输入的行跨度、尾部和紧凑逻辑宽度一起变化，不能把跨形状的时间差只归给 mask。

128×1024 下，8 warps 的三个进程 median 范围为 11.520–34.714 μs，明显比其他项宽。图中保留该范围，没有删掉慢进程；这组结果不足以支持稳定的配置优势。主形状重新运行紧凑 4 warps 得到 37.400 μs，与首次扫描接近，但只重跑一个配置不能重新证明它优于全部候选。1025 列的两配置随后在新进程中确认：4 warps 为 72.959 μs（72.639–73.039），8 warps 为 62.799 μs（61.479–62.959）。因此对这个已测形状保留紧凑宽度与 8 warps，见 图 10.11。

这一轮的可迁移方法是：先选覆盖一行的自然宽度，再固定它比较 warps，最后在固定 warps 下验证加宽是否值得。到了新形状，先重新算 B 和 program 数，再决定是否沿用原配置。

10.6 比较结果与保留方案

两条路线围绕同一公式，却直接控制不同的工作：

问题	HIP cooperative	HIP fused	Triton
行内分工	一 block 一行	相同行分工	一 program 一行
一次调用的主体 kernel 数	3	1	1
指数如何用于输出	全局保存后读取	再次计算	复用 program 内逻辑值
max、sum 归约	LDS 树	LDS 树与正确的复用屏障	tl.max、tl.sum
本轮主要选择	作为融合前的路线基线	中间交接与重算是否划算	warps 和逻辑宽度分别是否值得改变

配置	完整 median（μs）	三进程范围（μs）	相对本路线基线的速度比
hip-cooperative-b256	149.296	149.156–149.956	1.00×
hip-fused-b256	119.717	119.717–119.737	1.25×
triton-b4	37.599	37.199–37.599	1.00×
triton-b8	41.439	41.399–41.639	0.91×
triton-2b4	51.879	51.879–51.999	0.72×
triton-2b8	56.279	56.199–56.318	0.67×

上表为主输入 4096×1024。速度比用各自路线基线的时间除以当前时间；HIP 以 cooperative 为分母，Triton 以 (B,4) 为分母，不据此排列语言优劣。

图 10.11　独立确认使用 RX 9070 XT、ROCm 10.0 与完整输出时间。上面比较主形状的 HIP 融合，下面比较 1025 列的 Triton warps；每项仍为三进程各 50 次计时，两组分别统计。

HIP 主形状的融合收益在独立确认中保持，因而保留 fused；Triton 在 1024 列时保留紧凑宽度与 4 warps，1025 列则需要重新比较 warps。后一个案例说明，调参不是给算子贴上一个永远适用的配置标签。

判断融合时，把减少的中间读写、调用次数，与增加的重读和重算一起考虑。判断参数时，先确认是否只改了计划中的选择，再看有效工作和编译配置怎样变化。配置更大、逻辑访问更少，都只是待验证的线索。

10.7 运行与诊断

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

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

主阶段固定 4096×1024，先检查 28 项小边界，再测七项配置。 每项三个独立进程，各预热 10 次、计时 50 次，轮换候选顺序；先检查正确性，再汇总性能。main 不采集 profiler。图中心为三个进程中位数的中位数，误差线为它们的最小值到最大值，不是置信区间。

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

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

补测 4096×1023、4096×1025、128×1024，均保留 cooperative/fused 与四项 Triton。这一步继续使用本轮冻结的源码和配置，并检查实际环境。两条绘图命令都显式读取你的本轮 summary，不会用教程发布的成绩替缺失结果补值。

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

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

默认保留 cooperative、fused、Triton 基线与主形状最快候选，重复项只运行一次。确认数据保存在自己的 summary/summary.csv、summary/process-summary.csv 与 confirmation.json 中，不自动并入父实验。先比较同形状的新旧进程范围，再决定是否保留候选；不要复制教程中的确认文件到新目录拼图。

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

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

要检查 1025 列的 warps 排名是否保持，先完成 validation，再另建确认目录：

python chapter10/verify_rounds.py --frozen-run "$RUN/rounds" \
  --output "$RUN/confirmation-1025" \
  --shape 4096x1025 --configs triton-b4 triton-b8

此确认只回答新形状内的 4/8 warps 选择。主形状的 trace 不能替 1025 列解释原因；少量行的大波动也应保留，不能用最快一次替换。

10.8 练习：预测成本并验证

1. 稳定性与补齐：分别对 [0,0,0]、[1000,1000,1000] 手算稳定 Softmax，再用宽度 4 处理 [-1002,-1001,-1000]。如果无效位置填 0，它会怎样影响 max 和分母？
2. 构造关系测试：用同一基础行生成 x、x+1000、x−1000，检查平移前后的输出，并记录非有限结果。与“不同的行各自加一个平移量”相比，这项测试多检查了什么？
3. HIP 修改：保持 cooperative 与 fused 的相同行分工，在一个新形状上比较完整时间。若准备保存指数以避免重算，先说明每个线程要保存几个值、怎样处理尾部，再检查实际资源与性能；不能仅凭少一次 expf 宣布更快。
4. Triton 修改：选择一个新的列宽，先算 B，分别比较 (B,4)/(B,8) 与固定 warps 的 B/2B。保留全部结果，在新进程中确认候选；说明选择是否依赖行数。

自检提示

两组三个相同输入都得到 [1/3,1/3,1/3]。最大值不同，但平移后都成为 [0,0,0]。

全负输入若用 0 填充，无效位置可能成为最大值，并向分母加入自己的指数。真实三项还可能下溢为 0；即使最后的 store 不写无效位置，中间计算已经被污染。

保存指数会减少重算，也增加需要保留的值。扩大逻辑宽度可能改变编译分工，但在原宽度已覆盖整行时没有增加有效输入。这些改动的取舍都需要回到完整结果与实测时间。

本章小结

Softmax 将两次行内归约与逐元素计算连在一起。减最大值保护指数范围，单位元和 mask 让尾部正确参与运算；复用 LDS 时还要确保所有线程读完旧值，再覆盖为新用途。

HIP 先保持三个阶段建立行内协作基线，再比较融合，因而能更清楚地观察中间交接与重算的净效果。Triton 分别比较 warps 和逻辑宽度，并检查列宽边界与行数。每条路线都保留自己的初版和本轮对照，用新测量决定保留方案。

下一章进入矩阵乘。复用将从“同一行的多个阶段”扩展为“多个输出共同使用同一块输入”。

延伸阅读

• Triton Fused Softmax 教程：稳定公式、mask 与行融合；本章采用一行一个 program。
• HIP Kernel Language：共享存储与 block 同步。
• 第 9 章：Reduction：树形归约、单位元和阶段边界。
• 第 6 章：用 rocprof 找到慢点：筛选目标记录和解释字段。