第8章 Element-Wise：逐元素算子

本章导读

在第 1 篇中，我们以向量加法为例跑通了基准测试与性能剖析的全流程。进入算子实战的第一站，我们依然聚焦于看似最简单的加法运算 C[i] = A[i] + B[i]，但这次的核心命题转变为：当我们手里有了一个能跑通的正确基线后，如何系统性地判断下一步值得改什么？改动之后又如何科学地决定该保留还是回退？
本章我们将分别沿着 HIP 与 Triton 两条路线展开优化探索。你将学习到如何通过受控实验逐一检验“相邻访存连续性”、“单线程循环工作量”与“向量化加载（float4）”对性能的真实影响；当遇到没有加速甚至反常变慢的尝试时，学会坦然记录并作出理性的工程决策。第 4–7 章建立的计时基准、Trace 分析与 Roofline 模型是本章的基石；文中的具体语法可按需随时查阅附录。

配套代码编程范式附录

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

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

你可以先阅读两条路线共通的数学定义与数据流动画，再根据技术背景选择对应的实现入口：

阅读路线与前置知识

阅读路线	推荐起点
第一次接触 GPU 算子编程	建议先阅读附录 D：编程范式，建立基本图景后回到 8.1 节
倾向底层控制的 HIP 路线	阅读 8.1–8.3 节公共基础，在实现节选择 HIP 标签，最后阅读跨路线总结
倾向现代分块抽象的 Triton 路线	阅读 8.1–8.3 节公共基础，在实现节选择 Triton 标签，最后阅读跨路线总结

8.1 逐元素算子的数据依赖

8.1.1 输出之间的独立性

先看这个算子的数据依赖：一个输出需要哪些输入？计算它时，是否必须等待其他输出？

向量加法（Vector Add）的数学定义是：

C_{i} = A_{i} + B_{i}, 0 \leq i < N

计算 C[3] 时，只需读取 A[3] 与 B[3] 并相加，不需要等待 C[0]、C[1] 或其他位置的结果。

因此，不同输出之间互不依赖。每个位置的计算都可以看作一项独立任务。

8.1 语义层 · Element-Wise位置 i 独立计算 C[i] = A[i] + B[i]

位置 0：只读自己的两个输入位置 0 只读取 A[0]=2 与 B[0]=7，得到 C[0]=9，其他位置不受影响。

算法过程示意 · 数字定位到该步完成状态，可拖动进度或播放；动画速度不代表 GPU 耗时。

图 8.1　Vector Add 的逐元素依赖：当前位置只读取对应的两个输入位置。

图 8.1 每一步只高亮同一列的 A[i]、B[i] 与 C[i]，用来说明一个输出依赖哪些输入。实际程序可以把不同位置分给多个线程并行处理，也可以让一个线程处理多个位置。计算任务最终由 GPU 的计算单元（CU）执行，任务数量不等于 CU 数量；动画的逐列展示也不代表真实执行顺序。

现代深度学习中大量的核心算子都属于这种逐元素算子（Element-Wise Operators）：

算子	单个输出的表达式	`Y[i]` 依赖什么
Add（加法）	`Y[i] = A[i] + B[i]`	仅依赖同索引的 `A[i]` 与 `B[i]`
ReLU（激活）	`Y[i] = max(X[i], 0)`	仅依赖同索引的 `X[i]`
Scale & Bias（缩放平移）	`Y[i] = alpha * X[i] + beta`	仅依赖 `X[i]` 与全局常量标量
Clamp（截断）	`Y[i] = min(max(X[i], low), high)`	仅依赖 `X[i]` 与常量上下界

这些算子的共同点是：计算一个输出不需要其他输出的结果。程序可以通过索引公式与启动配置安排数据分工，再由硬件调度执行。同样的数学公式可以采用不同分工，后面的实验会比较这些选择对时间的影响。

8.1.2 HIP 与 Triton 的实现视角

虽然数学语义相同，但 HIP 与 Triton 采用了两种截然不同的抽象视角来把这项任务分工给硬件：

HIP 采用“线程主导”视角（Thread-centric）：代码描述的是单个线程的行为——“我是第几号线程，我该计算哪一个下标 $i$ ，我把对应元素读进来、相加、并写回”；
Triton 采用“分块主导”视角（Block/Tile-centric）：代码描述的是一组连续数据的集合行为——“我是一个计算实例（Program），我圈定一段逻辑索引范围，整体加载这组数据张量、执行向量加法并整体写回”。

如果你希望进一步厘清 Thread 与 Program、Block 与 Tile 的概念映射，请随时参考附录 D：HIP 与 Triton 的编程范式。

8.1.3 与归约算子的区别

为了更深刻地理解逐元素算子的特性，我们可以对比另一种截然相反的数据流动模式——归约（Reduction）。

求和归约的单个输出需要汇总全部输入（ $Y = \sum X_{i}$ ）；Softmax 的每个输出也需要用到整行的指数和。并行实现时，我们需要安排多个线程怎样合并这些数据。这是第 9 章 Reduction与第 10 章 Normalization要解决的问题。

因此请记住：“输入输出形状一致”并不等于逐元素算子。逐元素算子的本质边界在于：任意输出位置 $i$ 仅由对应输入位置 $i$ 的数值决定，不发生任何跨维度的跨行跨列跨位置聚合。

8.2 固定数学语义与正确性标准

8.2.1 四元素示例

我们用一个仅包含 4 个元素的小数组作为微型示范：

text

A = [2, -1, 4, 3]
B = [7,  3, -1, 2]

按照逐位置相加的规则：

text

C[0] = 2  + 7  = 9
C[1] = -1 + 3  = 2
C[2] = 4  + -1 = 3
C[3] = 3  + 2  = 5

C = [9, 2, 3, 5]

这个微型算例涵盖了该算子的全部数学逻辑。当我们将数组规模扩大到几千万个元素时，底层数学运算并未发生任何变化，改变的仅仅是我们如何在物理芯片上组织并发访存。

8.2.2 参考实现与检查条件

在优化过程中，“快但算错”没有任何意义。本章配套测试体系为 HIP 与 Triton 制定了完全统一的正确性红线：

项目	固定方式
数学语义	`C[i] = A[i] + B[i]`
数据类型	32 位浮点数（FP32）
输入	用同一确定性整数公式生成，再转成 FP32
参考结果	CPU 逐元素执行同一个 FP32 加法
边界	覆盖小于 wavefront、刚好等于 block、比 block 多 1、不能被 4 整除等长度
正确性	正式计时前检查一次，计时后再检查一次
计时范围	GPU event 包围一次 kernel 提交，不含分配、输入生成与 Host-to-Device 拷贝

为什么要测试非整除长度？例如，把 $N = 1027$ 个元素按每块 256 个元素划分，最后一块只剩 3 个有效元素。若采用“一线程处理一个元素、每 block 256 个线程”的初版 HIP 映射，最后一个 block 中就只有前 3 个线程需要读写数据，其余线程应跳过读写。

这里的线程数量取决于映射方式。后面的 grid-stride 循环与四元素分组需要按各自的索引检查尾部；Triton 的 mask 则标记逻辑元素位置，3 个有效位置不能直接理解成 3 个线程。无论怎样分工，都需要阻止越界读写。

Triton 测试脚本用 PyTorch 在 CPU 上计算参考结果，host_a 与 host_b 是两组输入：

python

reference = host_a + host_b

两条路线都逐个核对输出，要求结果有限、绝对误差不超过 1e-6。通过这些检查后，才比较运行时间。

8.2.3 非整除长度的边界

处理尾部边界时，两种路线展现了各自的表达习惯：

HIP 初版用 if (index < size) 保护单个元素的读写；使用循环的版本，则在每次迭代时检查下标；
Triton 用张量掩码（Mask）mask = offsets < size 标记有效位置，并将它传给加载和写回操作。

无论采用哪种写法，都必须同时覆盖读操作与写操作的边界防护。

8.3 建立成本模型和瓶颈假设

8.3.1 逻辑计算量与数据量

针对 FP32 单精度下的单元素向量加法，我们来数一数最理想情况下的逻辑开销：

动作	数量	逻辑字节
读取 `A[i]`	1 个 FP32	4 Byte
读取 `B[i]`	1 个 FP32	4 Byte
执行加法	1 次浮点加法	1 FLOP
写回 `C[i]`	1 个 FP32	4 Byte
合计		12 Byte + 1 FLOP

由此推导出其算法理论算术强度：

A I_{l o g i c a l} = \frac{1 FLOP}{12 Byte} \approx 0.0833 FLOP/Byte

必须明确：这是由算法逻辑读写量推导得出的逻辑算术强度，而不是硬件总线上测得的物理传输比例。在真实芯片中，由于片上缓存复用、显存事务打包或非对齐造成的额外搬运，各存储层级实际承受的数据流量都会与算法逻辑字节数存在差异。

基于这一逻辑算术强度，我们可以确立本章优化的首要指导假设：

向量加法每读写 12 字节逻辑数据仅执行 1 次浮点计算，逻辑算术强度极低；在大规模数据输入下，算子更容易受限于显存数据供给，而非浮点计算能力。

我们在第 7 章已经系统学习了 Roofline 的分析原则。这一假设提示我们：优化工作应优先从内存访存排布、合并事务与总线利用率入手排查，同时注意控制核函数启动、线程分工与循环展开等结构性开销。

8.3.2 有效带宽怎样算

本章所有版本的性能横向对比，均统一按照如下标准折算算法有效带宽：

B W_{e f f e c t i v e} = \frac{3 \times N \times 4 Byte}{t}

其中 t 是包围一次 kernel 提交的 GPU event 时间。这个指标适合在相同语义、相同 shape、相同计时范围下比较版本，但它不等于内存控制器实际传输了多少字节。物理流量必须由可用的硬件计数器或更进一步的分析支持。

8.4 实现向量加法

一个有用的实验从问题开始。例如，相邻线程的地址分散后会怎样；让更少的 block 各自做更多工作是否划算；让一个 Triton program 处理更多元素是否值得。假设可以来自前面的数据依赖和成本模型，再由测量检验。

先区分三件事，避免把“测到了一个数字”误当作解释：

方法	回答的问题	何时执行
正确性检查	输出仍符合加法语义吗？尾部有没有漏写或越界？	每次改动后，先检查再计时
Benchmark（性能测量）	相同条件下更快吗？收益超过运行波动了吗？	每轮改动都做，与相关对照一起比较
Profiling（性能分析）	哪项额外证据能帮助解释现象或排除一种猜测？	有具体问题且工具能回答时再做

如果已经决定撤回一项没有稳定收益的改动，就可以结束这一轮。只有还想回答“为何没有收益”或核对某项执行细节时，才需要进一步分析。下面两条路线各保留自然基线；编号用于区分实验，不代表逐级提速。

准备环境。 在 GPU 机器上，从仓库根目录进入本篇；自带的 uv.lock 用于锁定依赖，不需要删除：

bash

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

首次使用先按环境准备完成驱动、uv 和系统编译依赖。后文命令均在 code/part2-kernels/ 执行。

新建一轮实验。 在当前目录执行以下命令，保存打印出来的路径：

bash

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

每次新建都会得到不同目录，保留旧结果。这里的 RUN 只在当前终端有效。若换了终端，从仓库根目录重新执行环境准备，然后输入之前保存的目录：

bash

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

在提示处粘贴目录路径，不包含 RUN=，再按回车。只有看到确认信息才继续；目录不存在时先核对路径。继续原实验时不重建目录、不重复已有测量；需要重测或修改源码时，新建一轮并重新编译 HIP 程序。

命令开头会检查 RUN 是否已设置、目录是否存在，再进行写入。每条 benchmark 命令保存一个进程的 50 个样本。test ! -e ... && ... 防止追加到旧 CSV；文件已存在时程序不会运行。图中使用三个独立进程的结果，不能仅凭这里的一次运行挑选赢家。手工继续测第 2、3 个进程时，要同时将命令中的 benchmark-p1.samples.csv 和 --process 1 改为对应的 p2 / 2、p3 / 3；每一轮都依次运行本组全部候选，并轮换起始配置，保留所有结果。完成单项操作后，可用 8.8 的批量入口自动保存三进程比较与图。

固定比较条件。 本章数据来自 RX 9070 XT（gfx1201）、ROCm SDK 10.0.0、原生 Ubuntu 24.04；HIP 组件为 7.15.26333，Triton 为 3.8.0。

项目	所有逐轮图采用的条件
输入	`N=16,777,216`，FP32，确定性输入，`seed=20260920`
正确性	输出先填 NaN；计时前后检查所有结果有限、绝对误差不超过 `1e-6`
计时	预热 10 次、测量 50 次；分配、输入生成、拷贝、校验及 JIT 首次编译在计时之外
重复与图示	每配置 3 个独立进程；柱长为三个进程 median 的中位数，误差线为这三个 median 的最小—最大范围，不是置信区间
执行条件	轮换配置顺序；重复使用数组、不主动清缓存；保留桌面及已有后台上下文，不假定 GPU 独占

以下边界输出来自同一份 kernel 源码的独立正确性检查；它们与主输入的三进程计时分开保存。通过这些已选长度不能证明所有输入均正确，修改数学表达式或访问规则后仍需重新检查。

先看输出中的 precheck、postcheck、correct 是否全为 OK，再核对输入与配置，最后看 median_ms。effective_bandwidth_gbs 由逻辑字节数换算，不是硬件流量测量。计时口径沿用第 5 章，工具基础见第 6 章。

8.4.1 HIP 基线：一个线程处理一个元素

本章的起点是自然的一线程一元素写法。相邻线程访问相邻元素，数组加法又没有跨元素依赖，所以不需要先加入 LDS 或线程间同步。

cpp

__global__ void vector_add_v0(const float* __restrict__ input_a,
                              const float* __restrict__ input_b,
                              float* __restrict__ output,
                              std::size_t size) {
    const std::size_t index =
        static_cast<std::size_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    if (index < size) {
        output[index] = input_a[index] + input_b[index];
    }
}

N=16,777,216、每 block 256 个线程时，共需 65,536 个 block。最后一个 block 仍保留越界判断，以便同一份代码处理非整除长度。完整程序先准备输入与输出，对照 CPU 参考结果，再预热和计时。先看 vector_add_v0()，再看负责提交的 launch() 与组织测量的 run_version()。

完整源码：code/part2-kernels/chapter8/vector_add_hip.hip

cpp

// Chapter 8: Element-Wise Vector Add, HIP route.
//
// Performance results from this program are valid for the book only after it
// runs on the project baseline: Radeon RX 9070 XT, ROCm 10.0, native Ubuntu
// 24.04. The kernels deliberately keep the same c[i] = a[i] + b[i] semantics.

#include <hip/hip_runtime.h>

#include <algorithm>
#include <cmath>
#include <cctype>
#include <cstddef>
#include <cstdio>
#include <cstdlib>
#include <exception>
#include <limits>
#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 int kWaveSize = 32;
constexpr int kRoundsPerLane = 32;

__global__ void vector_add_v0(const float* __restrict__ input_a,
                              const float* __restrict__ input_b,
                              float* __restrict__ output,
                              std::size_t size) {
    const std::size_t index =
        static_cast<std::size_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    if (index < size) {
        output[index] = input_a[index] + input_b[index];
    }
}

// Each wave handles the same 32 x 32-element tile in both v1 kernels. Only
// the order in which lanes visit that tile changes, so grid size, loop count,
// arithmetic, and total logical bytes remain fixed.
__global__ void vector_add_v1_contiguous(const float* __restrict__ input_a,
                                         const float* __restrict__ input_b,
                                         float* __restrict__ output,
                                         std::size_t size) {
    const int waves_per_block = blockDim.x / kWaveSize;
    const int wave_in_block = threadIdx.x / kWaveSize;
    const int lane = threadIdx.x % kWaveSize;
    const std::size_t wave =
        static_cast<std::size_t>(blockIdx.x) * waves_per_block + wave_in_block;
    const std::size_t base = wave * kWaveSize * kRoundsPerLane;

#pragma unroll 1
    for (int round = 0; round < kRoundsPerLane; ++round) {
        const std::size_t index =
            base + static_cast<std::size_t>(round) * kWaveSize + lane;
        if (index < size) {
            output[index] = input_a[index] + input_b[index];
        }
    }
}

__global__ void vector_add_v1_strided(const float* __restrict__ input_a,
                                      const float* __restrict__ input_b,
                                      float* __restrict__ output,
                                      std::size_t size) {
    const int waves_per_block = blockDim.x / kWaveSize;
    const int wave_in_block = threadIdx.x / kWaveSize;
    const int lane = threadIdx.x % kWaveSize;
    const std::size_t wave =
        static_cast<std::size_t>(blockIdx.x) * waves_per_block + wave_in_block;
    const std::size_t base = wave * kWaveSize * kRoundsPerLane;

#pragma unroll 1
    for (int round = 0; round < kRoundsPerLane; ++round) {
        const std::size_t index =
            base + static_cast<std::size_t>(lane) * kRoundsPerLane + round;
        if (index < size) {
            output[index] = input_a[index] + input_b[index];
        }
    }
}

__global__ void vector_add_v2_grid_stride(const float* __restrict__ input_a,
                                          const float* __restrict__ input_b,
                                          float* __restrict__ output,
                                          std::size_t size) {
    const std::size_t thread =
        static_cast<std::size_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    const std::size_t grid_stride =
        static_cast<std::size_t>(gridDim.x) * blockDim.x;

    for (std::size_t index = thread; index < size; index += grid_stride) {
        output[index] = input_a[index] + input_b[index];
    }
}

__global__ void vector_add_v3_float4(const float* __restrict__ input_a,
                                     const float* __restrict__ input_b,
                                     float* __restrict__ output,
                                     std::size_t size) {
    const auto* input_a4 = reinterpret_cast<const float4*>(input_a);
    const auto* input_b4 = reinterpret_cast<const float4*>(input_b);
    auto* output4 = reinterpret_cast<float4*>(output);

    const std::size_t thread =
        static_cast<std::size_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    const std::size_t grid_stride =
        static_cast<std::size_t>(gridDim.x) * blockDim.x;
    const std::size_t vector_count = size / 4;

    for (std::size_t vector_index = thread;
         vector_index < vector_count;
         vector_index += grid_stride) {
        const float4 a = input_a4[vector_index];
        const float4 b = input_b4[vector_index];
        output4[vector_index] =
            make_float4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w);
    }

    // hipMalloc returns suitably aligned base pointers. The scalar loop keeps
    // the final 0-3 values correct when size is not divisible by four.
    for (std::size_t index = vector_count * 4 + thread;
         index < size;
         index += grid_stride) {
        output[index] = input_a[index] + input_b[index];
    }
}

struct Args {
    std::string version = "all";
    std::size_t size = 16ULL * 1024 * 1024;
    int block = 256;
    int grid = 0;
    int warmup = 10;
    int repeat = 50;
    std::size_t seed = 20260716;
    std::string samples;
    std::string config_id;
    int process = 1;
};

struct Version {
    std::string name;
    enum class Kind {
        v0,
        v1_contiguous,
        v1_strided,
        v2,
        v3,
    } kind;
};

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

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

[[noreturn]] void usage_error(const std::string& message) {
    std::fprintf(
        stderr,
        "%s\nUsage: vector_add_hip [--version all|v0|v1|v1-contiguous|"
        "v1-strided|v2|v3] [--size N] [--block N] [--grid 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 == "--size") {
                args.size = std::stoull(next());
            } else if (option == "--block") {
                args.block = std::stoi(next());
            } else if (option == "--grid") {
                args.grid = 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.size == 0) {
        usage_error("--size must be greater than zero");
    }
    if (args.block <= 0 || args.block > 1024 || args.block % kWaveSize != 0) {
        usage_error("--block must be a positive multiple of 32 and at most 1024");
    }
    if (args.grid < 0 || args.warmup < 0 || args.repeat <= 0) {
        usage_error("--grid/--warmup must be non-negative and --repeat positive");
    }
    if (args.process <= 0 || !std::all_of(args.config_id.begin(), args.config_id.end(),
            [](unsigned char c) { return std::isalnum(c) || c == '-' || c == '_' || c == '.'; })) {
        usage_error("--process must be positive; --config-id uses letters, digits, '-', '_' or '.'");
    }
    return args;
}

std::vector<Version> selected_versions(const std::string& selection) {
    const Version v0{"hip-v0", Version::Kind::v0};
    const Version v1_contiguous{
        "hip-v1-contiguous", Version::Kind::v1_contiguous};
    const Version v1_strided{"hip-v1-strided", Version::Kind::v1_strided};
    const Version v2{"hip-v2", Version::Kind::v2};
    const Version v3{"hip-v3", Version::Kind::v3};

    if (selection == "all") {
        return {v0, v1_contiguous, v1_strided, v2, v3};
    }
    if (selection == "v0") {
        return {v0};
    }
    if (selection == "v1") {
        return {v1_contiguous, v1_strided};
    }
    if (selection == "v1-contiguous") {
        return {v1_contiguous};
    }
    if (selection == "v1-strided") {
        return {v1_strided};
    }
    if (selection == "v2") {
        return {v2};
    }
    if (selection == "v3") {
        return {v3};
    }
    usage_error("unknown version: " + selection);
}

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 grid) {
    if (grid == 0 || grid > std::numeric_limits<unsigned int>::max()) {
        usage_error("calculated grid is outside the supported range");
    }
    return static_cast<unsigned int>(grid);
}

unsigned int grid_for(const Version& version,
                      const Args& args,
                      const hipDeviceProp_t& properties) {
    if (version.kind == Version::Kind::v0) {
        return checked_grid(ceil_div(args.size, args.block));
    }
    if (version.kind == Version::Kind::v1_contiguous ||
        version.kind == Version::Kind::v1_strided) {
        return checked_grid(ceil_div(
            args.size,
            static_cast<std::size_t>(args.block) * kRoundsPerLane));
    }

    const std::size_t work_items = version.kind == Version::Kind::v3
        ? std::max<std::size_t>(1, args.size / 4)
        : args.size;
    const std::size_t full_grid = ceil_div(work_items, args.block);
    if (args.grid > 0) {
        return checked_grid(std::min<std::size_t>(args.grid, full_grid));
    }
    const std::size_t persistent_grid =
        std::max(1, properties.multiProcessorCount) * 8ULL;
    return checked_grid(std::min(full_grid, persistent_grid));
}

void launch(const Version& version,
            unsigned int grid,
            int block,
            const float* input_a,
            const float* input_b,
            float* output,
            std::size_t size) {
    const dim3 grid_dim(grid);
    const dim3 block_dim(static_cast<unsigned int>(block));
    switch (version.kind) {
        case Version::Kind::v0:
            hipLaunchKernelGGL(vector_add_v0, grid_dim, block_dim, 0, 0,
                               input_a, input_b, output, size);
            break;
        case Version::Kind::v1_contiguous:
            hipLaunchKernelGGL(vector_add_v1_contiguous,
                               grid_dim, block_dim, 0, 0,
                               input_a, input_b, output, size);
            break;
        case Version::Kind::v1_strided:
            hipLaunchKernelGGL(vector_add_v1_strided,
                               grid_dim, block_dim, 0, 0,
                               input_a, input_b, output, size);
            break;
        case Version::Kind::v2:
            hipLaunchKernelGGL(vector_add_v2_grid_stride,
                               grid_dim, block_dim, 0, 0,
                               input_a, input_b, output, size);
            break;
        case Version::Kind::v3:
            hipLaunchKernelGGL(vector_add_v3_float4,
                               grid_dim, block_dim, 0, 0,
                               input_a, input_b, output, size);
            break;
    }
    HIP_CHECK(hipGetLastError());
}

void make_inputs(std::vector<float>& input_a,
                 std::vector<float>& input_b,
                 std::vector<float>& reference,
                 std::size_t seed) {
    for (std::size_t index = 0; index < input_a.size(); ++index) {
        const long long a_integer = static_cast<long long>(
            (index * 17ULL + seed) % 2048ULL) - 1024;
        const long long b_integer = static_cast<long long>(
            (index * 29ULL + seed * 3ULL) % 4096ULL) - 2048;
        input_a[index] = static_cast<float>(a_integer) / 1024.0F;
        input_b[index] = static_cast<float>(b_integer) / 2048.0F;
        reference[index] = input_a[index] + input_b[index];
    }
}

Validation validate(const float* device_output,
                    const std::vector<float>& reference,
                    std::vector<float>& host_output) {
    HIP_CHECK(hipMemcpy(host_output.data(), device_output,
                        host_output.size() * sizeof(float),
                        hipMemcpyDeviceToHost));

    double max_error = 0.0;
    std::size_t first_bad = reference.size();
    for (std::size_t index = 0; index < reference.size(); ++index) {
        const double error = std::abs(
            static_cast<double>(host_output[index]) - reference[index]);
        max_error = std::max(max_error, error);
        if ((!std::isfinite(host_output[index]) || error > 1.0e-6) &&
            first_bad == reference.size()) {
            first_bad = index;
        }
    }
    return {first_bad == reference.size(), max_error, first_bad};
}

Timing benchmark(const Version& version,
                 unsigned int grid,
                 const Args& args,
                 const float* input_a,
                 const float* input_b,
                 float* output) {
    for (int iteration = 0; iteration < args.warmup; ++iteration) {
        launch(version, grid, args.block, input_a, input_b, output, args.size);
    }
    HIP_CHECK(hipDeviceSynchronize());

    std::vector<hipEvent_t> starts(args.repeat);
    std::vector<hipEvent_t> ends(args.repeat);
    for (int iteration = 0; iteration < args.repeat; ++iteration) {
        HIP_CHECK(hipEventCreate(&starts[iteration]));
        HIP_CHECK(hipEventCreate(&ends[iteration]));
    }

    for (int iteration = 0; iteration < args.repeat; ++iteration) {
        HIP_CHECK(hipEventRecord(starts[iteration]));
        launch(version, grid, args.block, input_a, input_b, output, args.size);
        HIP_CHECK(hipEventRecord(ends[iteration]));
    }
    HIP_CHECK(hipDeviceSynchronize());

    std::vector<double> times_ms(args.repeat);
    for (int iteration = 0; iteration < args.repeat; ++iteration) {
        float elapsed_ms = 0.0F;
        HIP_CHECK(hipEventElapsedTime(
            &elapsed_ms, starts[iteration], ends[iteration]));
        times_ms[iteration] = elapsed_ms;
        HIP_CHECK(hipEventDestroy(starts[iteration]));
        HIP_CHECK(hipEventDestroy(ends[iteration]));
    }

    const std::vector<double> samples_ms = times_ms;
    std::sort(times_ms.begin(), times_ms.end());
    const std::size_t middle = times_ms.size() / 2;
    const double median = times_ms.size() % 2 == 0
        ? (times_ms[middle - 1] + times_ms[middle]) / 2.0
        : times_ms[middle];
    double sum = 0.0;
    for (const double value : times_ms) {
        sum += value;
    }
    return {times_ms.front(), median, sum / times_ms.size(), samples_ms};
}

void write_samples(const Version& version, const Args& args,
                   unsigned int grid, const Timing& timing) {
    if (args.samples.empty()) return;
    std::FILE* file = std::fopen(args.samples.c_str(), "a+");
    if (!file) usage_error("cannot open sample CSV: " + args.samples);
    std::fseek(file, 0, SEEK_END);
    if (std::ftell(file) == 0) {
        std::fprintf(file, "config_id,implementation,runtime,shape,dtype,block,grid,num_warps,warmup,repeat,seed,process,sample,ms\n");
    }
    const std::string id = args.config_id.empty() ? version.name : args.config_id;
    for (std::size_t i = 0; i < timing.samples_ms.size(); ++i) {
        std::fprintf(file, "%s,%s,hip,%zu,float32,%d,%u,0,%d,%d,%zu,%d,%zu,%.12g\n",
                     id.c_str(), version.name.c_str(), args.size, args.block, grid,
                     args.warmup, args.repeat, args.seed, args.process, i,
                     timing.samples_ms[i]);
    }
    if (std::ferror(file)) { std::fclose(file); usage_error("failed to write sample CSV"); }
    if (std::fclose(file) != 0) usage_error("failed to close sample CSV");
}

bool run_version(const Version& version,
                 const Args& args,
                 const hipDeviceProp_t& properties,
                 const float* device_a,
                 const float* device_b,
                 float* device_output,
                 const std::vector<float>& reference,
                 std::vector<float>& host_output) {
    const unsigned int grid = grid_for(version, args, properties);
    const std::size_t bytes = args.size * sizeof(float);

    HIP_CHECK(hipMemset(device_output, 0xFF, bytes));
    launch(version, grid, args.block,
           device_a, device_b, device_output, args.size);
    HIP_CHECK(hipDeviceSynchronize());
    const Validation precheck =
        validate(device_output, reference, host_output);
    if (!precheck.correct) {
        std::fprintf(
            stderr,
            "precheck failed for %s at index %zu: got=%g expected=%g\n",
            version.name.c_str(), precheck.first_bad_index,
            host_output[precheck.first_bad_index],
            reference[precheck.first_bad_index]);
        std::printf(
            "RESULT operator=vector-add implementation=%s runtime=hip "
            "shape=%zu dtype=float32 "
            "block=%d grid=%u warmup=%d repeat=%d seed=%zu timed=0 "
            "correct=FAIL precheck=FAIL postcheck=NA min_ms=NA "
            "median_ms=NA mean_ms=NA effective_bandwidth_gbs=NA "
            "max_abs_error=%.9g\n",
            version.name.c_str(), args.size, args.block, grid,
            args.warmup, args.repeat, args.seed,
            precheck.max_absolute_error);
        return false;
    }

    const Timing timing = benchmark(
        version, grid, args, device_a, device_b, device_output);
    const Validation postcheck =
        validate(device_output, reference, host_output);
    if (postcheck.correct) write_samples(version, args, grid, timing);
    const double logical_bytes =
        3.0 * static_cast<double>(args.size) * sizeof(float);
    const double effective_bandwidth_gbs =
        logical_bytes / (timing.median_ms * 1.0e-3) / 1.0e9;

    std::printf(
        "RESULT operator=vector-add implementation=%s runtime=hip "
        "shape=%zu dtype=float32 "
        "block=%d grid=%u warmup=%d repeat=%d seed=%zu timed=1 "
        "correct=%s precheck=OK postcheck=%s min_ms=%.6f "
        "median_ms=%.6f mean_ms=%.6f effective_bandwidth_gbs=%.6f "
        "max_abs_error=%.9g\n",
        version.name.c_str(), args.size, args.block, grid,
        args.warmup, args.repeat, args.seed,
        postcheck.correct ? "OK" : "FAIL",
        postcheck.correct ? "OK" : "FAIL",
        timing.minimum_ms, timing.median_ms, timing.mean_ms,
        effective_bandwidth_gbs, postcheck.max_absolute_error);
    return postcheck.correct;
}

}  // namespace

int main(int argc, char** argv) {
    const Args args = parse_args(argc, argv);
    const std::vector<Version> versions = selected_versions(args.version);

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

    const std::size_t bytes = args.size * sizeof(float);
    std::vector<float> host_a(args.size);
    std::vector<float> host_b(args.size);
    std::vector<float> reference(args.size);
    std::vector<float> host_output(args.size);
    make_inputs(host_a, host_b, reference, args.seed);

    float* device_a = nullptr;
    float* device_b = nullptr;
    float* device_output = nullptr;
    HIP_CHECK(hipMalloc(&device_a, bytes));
    HIP_CHECK(hipMalloc(&device_b, bytes));
    HIP_CHECK(hipMalloc(&device_output, bytes));
    HIP_CHECK(hipMemcpy(
        device_a, host_a.data(), bytes, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(
        device_b, host_b.data(), bytes, hipMemcpyHostToDevice));

    bool all_correct = true;
    for (const Version& version : versions) {
        all_correct = run_version(
            version, args, properties,
            device_a, device_b, device_output,
            reference, host_output) && all_correct;
    }

    HIP_CHECK(hipFree(device_a));
    HIP_CHECK(hipFree(device_b));
    HIP_CHECK(hipFree(device_output));
    return all_correct ? EXIT_SUCCESS : EXIT_FAILURE;
}

这份文件包含本路线所有候选。后面只需定位相应函数并改变运行参数，不必复制多份主机程序。先编译一次；已有可执行文件时直接复用：

bash

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

--offload-arch=gfx1201 选择本章显卡的编译目标。若改了源码，请新建 RUN 后重新编译，避免新源码与旧可执行文件混用。

运行与校验

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
mkdir -p "$RUN/configs/hip-v0" && \
test ! -e "$RUN/configs/hip-v0/benchmark-p1.samples.csv" && \
"$RUN/build/vector_add_hip" \
  --version v0 --block 256 \
  --size 16777216 --warmup 10 --repeat 50 --seed 20260920 \
  --samples "$RUN/configs/hip-v0/benchmark-p1.samples.csv" \
  --config-id hip-v0 --process 1

运行输出：hip-v0，第 1 个独立进程

text

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 multi_processor_count=32 wave_size=32
RESULT operator=vector-add implementation=hip-v0 runtime=hip shape=16777216 dtype=float32 block=256 grid=65536 warmup=10 repeat=50 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.348321 median_ms=0.349441 mean_ms=0.349336 effective_bandwidth_gbs=576.139024 max_abs_error=0

先核对 shape=16777216、block=256、grid=65536 和三个 OK。这次单进程的 median_ms=0.349441 是 50 次 event 时间的中位数。另两个独立进程分别为 0.333301、0.334861 ms；图中采用三者中位数 0.334861 ms，范围 0.333301–0.349441 ms。较慢的一次也保留，不能只取最快值。

基线分析：计时区间里实际执行了什么？ 这一处 trace 用来核对目标 kernel 的起止时间和实际启动规模，建立后续读报告的方法。它不负责证明“已经达到带宽极限”。

保持当前目录和 RUN，另启动一次带 profiler 的进程。分析时预热 5 次、记录 10 次；这次终端打印的 event 时间不进入 benchmark。inspect_trace.py 只用 Python 标准库读取 CSV，不运行 GPU。

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
mkdir "$RUN/configs/hip-v0/profile" && \
rocprofv3 --kernel-trace --hip-trace --stats \
  --output-format csv pftrace \
  --output-directory "$RUN/configs/hip-v0/profile" \
  --output-file hip-v0 -- \
  "$RUN/build/vector_add_hip" \
  --version v0 --block 256 \
  --size 16777216 --warmup 5 --repeat 10 --seed 20260920 \
  --config-id hip-v0 --process 1 && \
python chapter8/inspect_trace.py \
  "$RUN/configs/hip-v0/profile/hip-v0_kernel_trace.csv" \
  --kernel vector_add_v0 --skip 6 --take 10

性能分析原始输出：hip-v0（终端、生成记录与 CSV）

程序 stdout

text

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 multi_processor_count=32 wave_size=32
RESULT operator=vector-add implementation=hip-v0 runtime=hip shape=16777216 dtype=float32 block=256 grid=65536 warmup=5 repeat=10 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.335961 median_ms=0.337041 mean_ms=0.336981 effective_bandwidth_gbs=597.335629 max_abs_error=0

Profiler stderr：文件生成记录

text

W0920 08:17:17.555068 1401833 simple_timer.cpp:55] [rocprofv3] tool initialization ::     0.001453 sec
W0920 08:17:17.555337 1401833 tool.cpp:3348] HIP (compiler) version 7.15.0 initialized (instance=0)
W0920 08:17:17.555379 1401833 simple_timer.cpp:55] [rocprofv3] 'chapter8/results/walkthrough-20260920/build/vector_add_hip --version v0 --size 16777216 --block 256 --warmup 5 --repeat 10 --seed 20260920 --config-id hip-v0 --process 1' ::     0.000000 sec
W0920 08:17:17.555529 1401833 tool.cpp:3348] HIP (runtime) version 7.15.0 initialized (instance=0)
W0920 08:17:17.561246 1401833 tool.cpp:3348] HSA version 1.21.0 initialized (instance=0)
W0920 08:17:17.658814 1401833 queue_controller.cpp:170] [queue-intercept] device-memory ring-buffer requested but profiling requires system-memory InterceptQueue; falling back to system-memory ring for agent 100431044319616 (priority and CU-mask from the descriptor are not preserved)
W0920 08:17:17.767101 1401833 simple_timer.cpp:55] [rocprofv3] 'chapter8/results/walkthrough-20260920/build/vector_add_hip --version v0 --size 16777216 --block 256 --warmup 5 --repeat 10 --seed 20260920 --config-id hip-v0 --process 1' ::     0.211721 sec
E0920 08:17:17.796019 1401833 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/hip-v0/profile/hip-v0_kernel_stats.csv
E0920 08:17:17.796170 1401833 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/hip-v0/profile/hip-v0_kernel_trace.csv
E0920 08:17:17.813048 1401833 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/hip-v0/profile/hip-v0_hip_api_stats.csv
E0920 08:17:17.813188 1401833 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/hip-v0/profile/hip-v0_hip_api_trace.csv
E0920 08:17:17.818296 1401833 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/hip-v0/profile/hip-v0_agent_info.csv
E0920 08:17:17.818407 1401833 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/hip-v0/profile/hip-v0_domain_stats.csv
[153.869]       perfetto.cc:47314 Configured tracing session 1, #sources:1, duration:0 ms, #buffers:1, total buffer size:1048576 KB, total sessions:1, uid:0 session name: ""
E0920 08:17:17.846910 1401833 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/hip-v0/profile/hip-v0_results.pftrace
[153.888]       perfetto.cc:48889 Tracing session 1 ended, total sessions:0
W0920 08:17:17.849113 1401833 simple_timer.cpp:55] [rocprofv3] output generation ::     0.067913 sec
W0920 08:17:17.849147 1401833 simple_timer.cpp:55] [rocprofv3] tool finalization ::     0.067986 sec

原始 kernel trace CSV

text

"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",2,0,1401833,1,4,"__amd_rocclr_fillBufferUnAligned",14,908153729517605,908153729558685,0,0,8,0,128,256,1,1,8192,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401833,2,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",17,908153741205911,908153741882753,0,0,8,0,128,256,1,1,16777216,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401833,3,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",23,908153773225567,908153773685129,0,0,8,0,128,256,1,1,16777216,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401833,4,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",27,908153773689409,908153774155491,0,0,8,0,128,256,1,1,16777216,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401833,5,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",31,908153774160891,908153774733293,0,0,8,0,128,256,1,1,16777216,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401833,6,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",35,908153774747493,908153775111294,0,0,8,0,128,256,1,1,16777216,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401833,7,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",39,908153775116854,908153775481176,0,0,8,0,128,256,1,1,16777216,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401833,8,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",65,908153775686576,908153776012898,0,0,8,0,128,256,1,1,16777216,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401833,9,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",71,908153776029138,908153776355819,0,0,8,0,128,256,1,1,16777216,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401833,10,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",77,908153776371859,908153776698100,0,0,8,0,128,256,1,1,16777216,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401833,11,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",83,908153776714052,908153777040893,0,0,8,0,128,256,1,1,16777216,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401833,12,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",89,908153777056613,908153777382814,0,0,8,0,128,256,1,1,16777216,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401833,13,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",95,908153777398534,908153777725255,0,0,8,0,128,256,1,1,16777216,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401833,14,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",101,908153777741015,908153778067216,0,0,8,0,128,256,1,1,16777216,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401833,15,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",107,908153778083176,908153778409058,0,0,8,0,128,256,1,1,16777216,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401833,16,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",113,908153778424658,908153778751539,0,0,8,0,128,256,1,1,16777216,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401833,17,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",119,908153778767179,908153779092500,0,0,8,0,128,256,1,1,16777216,1,1

筛选后的分析结果：hip-v0

text

kernel_name="(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)"
csv_rows=17 target_rows=16 skipped_rows=6 kept_rows=10
sort=Start_Timestamp_ascending
first_start_ns=908153775686576 first_end_ns=908153776012898 first_duration_ns=326322
kernel_min_us=325.321000 kernel_median_us=326.281500 kernel_max_us=326.881000
grid_work_items_xyz=16777216x1x1 workgroup_size_xyz=256x1x1 workgroups_xyz=65536x1x1
workgroups_total=65536
vgpr_count=8 sgpr_count=128 lds_bytes=0 scratch_bytes=0

先找 vector_add_v0：原始 CSV 还包含填充输出的辅助 kernel，目标记录共有 16 条。去掉 1 次预检和 5 次预热，保留 10 次，这就是命令中 --skip 6 --take 10 的含义；整份 kernel_stats 的平均值包含预检与预热，不能直接作正式成绩。

再核对时间：第一条保留记录为 908153776012898 − 908153775686576 = 326322 ns，即 326.322 μs；10 次的中位数为 326.2815 μs。它来自独立采集，与 benchmark 的 event 时间不是同一次运行，不能相减来估算 CPU 开销。

最后核对启动规模：grid_work_items_xyz=16777216x1x1 除以 workgroup_size_xyz=256x1x1 得到 65,536 个 block。这里的 Grid X 是工作项数。vgpr_count=8、scratch_bytes=0 描述资源分配，不能单独证明动态占用率或带宽瓶颈。

这次核对确认了目标、规模和统计区间。下一项假设可以直接来自地址安排，无需先为每个候选再采一份相同报告。分析脚本全文和时间线读法见 8.5。

8.4.2 地址对照：只改变同轮访问顺序

基线已经采用连续访问；不能把故意改坏它再改回来称为新的优化。这里另设一对实验，检验“同轮地址排列是否值得关注”，再把得到的判断带回正常实现。

base 是这组线程所负责片段的起点，round 是当前循环轮次。真实代码让一个 wave32 处理 1,024 个元素：32 个 lane 各处理 32 个元素，共循环 32 轮。两版的 grid、block、循环轮数、加法次数和输入完全相同，仅下标排列不同：

cpp

// hip-v1-contiguous：同一轮，相邻 lane 读取相邻位置
index = base + round * 32 + lane;

// hip-v1-strided：同一轮，相邻 lane 的位置相隔 32 个元素
index = base + lane * 32 + round;

两种排法最终都会覆盖同一片输入，也产生同一份输出。区别发生在同一轮请求哪些地址。

8.4 HIP · 合并访存 · 8-lane 教学缩略同样 8 个 lane，地址排法决定触及几组

左右滑动画布，查看完整过程

同一批元素，两种下标排法同样 8 个 lane、同样 32 个元素。唯一变量：每轮内 lane 到下标的排列——左边 index = round×8+lane（连续），右边 index = lane×4+round（跨步）。

算法过程示意 · 数字定位到该步完成状态，可拖动进度或播放；动画速度不代表 GPU 耗时。

图 8.2　地址排列的缩略演示：8 个 lane 各处理 4 个元素。32B 分组只帮助观察地址是否集中，不是 RX 9070 XT 的实测事务计数。

图把真实代码的 32 个 lane × 32 轮缩成 8 个 lane × 4 轮，保留同轮连续或分散的差别。先在图 8.2 中看一轮：连续排列的地址集中在一起，跨步排列散布在多组。再播放全部四轮，确认两边最终没有遗漏或重复。真实性能要看独立计时：

两次运行分别调用同一文件的 vector_add_v1_contiguous() 与 vector_add_v1_strided()，由 --version 选择。完整主机程序见本路线源码。

运行与校验

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
mkdir -p "$RUN/configs/hip-v1-contiguous" && \
test ! -e "$RUN/configs/hip-v1-contiguous/benchmark-p1.samples.csv" && \
"$RUN/build/vector_add_hip" \
  --version v1-contiguous --block 256 \
  --size 16777216 --warmup 10 --repeat 50 --seed 20260920 \
  --samples "$RUN/configs/hip-v1-contiguous/benchmark-p1.samples.csv" \
  --config-id hip-v1-contiguous --process 1

运行输出：hip-v1-contiguous，第 1 个独立进程

text

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 multi_processor_count=32 wave_size=32
RESULT operator=vector-add implementation=hip-v1-contiguous runtime=hip shape=16777216 dtype=float32 block=256 grid=2048 warmup=10 repeat=50 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.366922 median_ms=0.370081 mean_ms=0.370749 effective_bandwidth_gbs=544.006097 max_abs_error=0

这是本轮的连续访问对照。输出应为 grid=2048、block=256，且前后检查通过；本次单进程中位数为 0.370081 ms。

运行与校验

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
mkdir -p "$RUN/configs/hip-v1-strided" && \
test ! -e "$RUN/configs/hip-v1-strided/benchmark-p1.samples.csv" && \
"$RUN/build/vector_add_hip" \
  --version v1-strided --block 256 \
  --size 16777216 --warmup 10 --repeat 50 --seed 20260920 \
  --samples "$RUN/configs/hip-v1-strided/benchmark-p1.samples.csv" \
  --config-id hip-v1-strided --process 1

运行输出：hip-v1-strided，第 1 个独立进程

text

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 multi_processor_count=32 wave_size=32
RESULT operator=vector-add implementation=hip-v1-strided runtime=hip shape=16777216 dtype=float32 block=256 grid=2048 warmup=10 repeat=50 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=2.427285 median_ms=2.538206 mean_ms=2.540035 effective_bandwidth_gbs=79.318476 max_abs_error=0

跨步版保持同样的输入、grid、block 与循环工作量，输出仍然正确；本次单进程中位数为 2.538206 ms。接下来用三进程结果确认差异是否只是偶然波动。

相同工作量下连续访问与跨步访问的时间比较

图 8.3　只改变同轮地址排列的 HIP 对照；柱长为三个独立进程 median 的中位数，误差线为进程范围。

连续版为 0.369681 ms，跨步版为 2.538206 ms，后者用时约为前者的 6.87 倍，两者的进程范围明显分离。

受控对照支持“同轮地址排列显著影响当前实现的时间”。它同时改变 A/B 的读取与 C 的写入排列，尚未把时间分摊给读、写或具体缓存行为；图中的地址分组也不是实测内存事务。基线 v0 本来就连续访问，因此这不是把 v0 加速了 6.87 倍。

本轮的决定是保留同轮相邻访问。再采一份仅重复“跨步版更慢”的 trace，不会改变这个决定。若要继续量化物理流量或缓存命中率，需要有效的硬件计数器；本机本次没有取得相应有效数据，边界见 8.5.3。

8.4.3 Grid-Stride：减少 block 数是否有收益

基线启动了 65,536 个 block，每个线程只做一次加法。一个可以验证的想法是：让线程循环处理多个位置，能否用更少的 block 完成同样的工作？

hip-v2 用整个 grid 的线程数作为步长：

cpp

const std::size_t thread =
    static_cast<std::size_t>(blockIdx.x) * blockDim.x + threadIdx.x;
const std::size_t grid_stride =
    static_cast<std::size_t>(gridDim.x) * blockDim.x;

for (std::size_t index = thread; index < size; index += grid_stride) {
    output[index] = input_a[index] + input_b[index];
}

假设整个 grid 有 1,024 个线程，线程 0 处理 0、1024、2048…，线程 1 处理 1、1025、2049…。单个线程跨着读，同一轮的相邻线程仍然连续读，与上一轮的跨步反例不同。

不能只把 v0 与一个缩小了 grid 的 v2 相比，就把差异都归给 block 数。因此分两步观察：先让 v2 同样使用 65,536 个 block，再固定这份循环代码，只把 grid 改为 2,048、256。输入、block size、计时方式保持不变。

本轮重点阅读完整程序中的 vector_add_v2_grid_stride() 与 grid_for()。三次使用同一份循环代码，分别给出明确的 grid 参数。

保持 65,536 个 block。

运行与校验

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
mkdir -p "$RUN/configs/hip-v2-g65536" && \
test ! -e "$RUN/configs/hip-v2-g65536/benchmark-p1.samples.csv" && \
"$RUN/build/vector_add_hip" \
  --version v2 --block 256 --grid 65536 \
  --size 16777216 --warmup 10 --repeat 50 --seed 20260920 \
  --samples "$RUN/configs/hip-v2-g65536/benchmark-p1.samples.csv" \
  --config-id hip-v2-g65536 --process 1

运行输出：hip-v2-g65536，第 1 个独立进程

text

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 multi_processor_count=32 wave_size=32
RESULT operator=vector-add implementation=hip-v2 runtime=hip shape=16777216 dtype=float32 block=256 grid=65536 warmup=10 repeat=50 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.354801 median_ms=0.355121 mean_ms=0.355214 effective_bandwidth_gbs=566.923949 max_abs_error=0

先核对 grid=65536 与 v0 相同。此时检验的是换成循环结构的效果，不能把它和后面的 grid 缩小混为一项改动。本次单进程中位数为 0.355121 ms。

减少到 2,048 个 block。

运行与校验

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
mkdir -p "$RUN/configs/hip-v2-g2048" && \
test ! -e "$RUN/configs/hip-v2-g2048/benchmark-p1.samples.csv" && \
"$RUN/build/vector_add_hip" \
  --version v2 --block 256 --grid 2048 \
  --size 16777216 --warmup 10 --repeat 50 --seed 20260920 \
  --samples "$RUN/configs/hip-v2-g2048/benchmark-p1.samples.csv" \
  --config-id hip-v2-g2048 --process 1

运行输出：hip-v2-g2048，第 1 个独立进程

text

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 multi_processor_count=32 wave_size=32
RESULT operator=vector-add implementation=hip-v2 runtime=hip shape=16777216 dtype=float32 block=256 grid=2048 warmup=10 repeat=50 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.359441 median_ms=0.361721 mean_ms=0.361690 effective_bandwidth_gbs=556.579759 max_abs_error=0

输出中的 grid=2048 表明配置已经生效。平均每线程处理 32 个元素，总加法数量没有减少；本次单进程中位数为 0.361721 ms。

减少到 256 个 block。

运行与校验

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
mkdir -p "$RUN/configs/hip-v2-g256" && \
test ! -e "$RUN/configs/hip-v2-g256/benchmark-p1.samples.csv" && \
"$RUN/build/vector_add_hip" \
  --version v2 --block 256 --grid 256 \
  --size 16777216 --warmup 10 --repeat 50 --seed 20260920 \
  --samples "$RUN/configs/hip-v2-g256/benchmark-p1.samples.csv" \
  --config-id hip-v2-g256 --process 1

运行输出：hip-v2-g256，第 1 个独立进程

text

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 multi_processor_count=32 wave_size=32
RESULT operator=vector-add implementation=hip-v2 runtime=hip shape=16777216 dtype=float32 block=256 grid=256 warmup=10 repeat=50 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.342720 median_ms=0.343640 mean_ms=0.343801 effective_bandwidth_gbs=585.864836 max_abs_error=0

此时平均每线程处理 256 个元素。确认 correct=OK 后，记录本次单进程中位数 0.343640 ms；不要因为这一项较小，就跳过重复测量。

固定 grid-stride kernel 后比较三种 grid 大小，并以原始基线作为参照

图 8.4　v0 是原始参照；v2 的三个配置使用同一份 kernel。缩小 grid 会同时增加每个线程的工作量，图中实际比较的是这项分工选择。

三个 v2 配置分别为 0.350301、0.357582、0.343640 ms，本轮都没有显示稳定优于 v0 的收益。减少 block 数会同时增加每个线程的循环量；预期中的管理成本下降，未必能抵消其他代价。

本轮保留 v0。这个决定只需要公平测量，不需要给三档 grid 分别再做完整 trace。若后续要判断是否启动了预期数量的 block，可针对那一档检查 Grid / Workgroup；若要解释运行时占用率或等待，还需要对应的有效证据，不能用 block 数替代它们。

8.4.4 `float4`：一次表达四个元素是否更快

上一轮的 grid=256 中位数高于基线，进程范围仍有重叠。每个线程要循环多次，按四个元素一组读写，能否弥补一部分开销？为了检验这个想法，我们继续固定 grid=256、block=256，比较标量循环 v2 与向量类型 v3；图中同时保留原始 v0。这里没有把 256 当作已经选出的最优 grid。

本例从 hipMalloc 返回的对齐基地址开始；任意 input + 1 偏移不再保证适合强转为 float4*，复用这段代码时要先保证地址对齐。

float4 把四个 FP32 组合成一个 16 Byte 类型。v3 按四元素组读写，并为剩余的 0–3 个元素保留标量路径：

8.4 HIP · float4 向量化四个 FP32 组成 float4，尾部单独处理

左右滑动画布，查看完整过程

先把 A 看成 18 个独立元素N=18 个 FP32，每个占 4 Byte——这是 v0/v1/v2 标量世界的做法。

算法过程示意 · 数字定位到该步完成状态，可拖动进度或播放；动画速度不代表 GPU 耗时。

图 8.5　N=18 的源码分组：前 16 个元素组成四组，最后两个元素单独处理。图中只展开输入 A，B 和 C 使用相同分组。

cpp

const std::size_t vector_count = size / 4;
for (std::size_t vector_index = thread;
     vector_index < vector_count;
     vector_index += grid_stride) {
    const float4 a = input_a4[vector_index];
    const float4 b = input_b4[vector_index];
    output4[vector_index] =
        make_float4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w);
}
for (std::size_t index = vector_count * 4 + thread;
     index < size;
     index += grid_stride) {
    output[index] = input_a[index] + input_b[index];
}

本轮的假设是向量类型可能改变生成的访存指令及循环工作量。源码写出 float4 本身不能证明事务减少；先通过包含尾部的正确性检查，再判断时间：

完整实现仍在HIP 源文件的 vector_add_v3_float4()，运行参数使用 --version v3。

先检查新增的尾部路径：N=2 只走标量尾部，N=6 同时包含一组 float4 和两个尾部元素，N=1026 检查较长向量部分后的两个元素。这里预热为 0、计时重复 1 次，程序仍会先做一次预检；目的是核对答案。

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
"$RUN/build/vector_add_hip" --version v3 --size 2 --block 256 --warmup 0 --repeat 1 --seed 20260920 --grid 256 && \
"$RUN/build/vector_add_hip" --version v3 --size 6 --block 256 --warmup 0 --repeat 1 --seed 20260920 --grid 256 && \
"$RUN/build/vector_add_hip" --version v3 --size 1026 --block 256 --warmup 0 --repeat 1 --seed 20260920 --grid 256

边界检查原始输出：float4 尾部 @ RX 9070 XT / ROCm 10.0

boundary-float4-n2.log

text

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 multi_processor_count=32 wave_size=32
RESULT operator=vector-add implementation=hip-v3 runtime=hip shape=2 dtype=float32 block=256 grid=1 warmup=0 repeat=1 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.011160 median_ms=0.011160 mean_ms=0.011160 effective_bandwidth_gbs=0.002151 max_abs_error=0

boundary-float4-n6.log

text

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 multi_processor_count=32 wave_size=32
RESULT operator=vector-add implementation=hip-v3 runtime=hip shape=6 dtype=float32 block=256 grid=1 warmup=0 repeat=1 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.011200 median_ms=0.011200 mean_ms=0.011200 effective_bandwidth_gbs=0.006429 max_abs_error=0

boundary-float4-n1026.log

text

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 multi_processor_count=32 wave_size=32
RESULT operator=vector-add implementation=hip-v3 runtime=hip shape=1026 dtype=float32 block=256 grid=1 warmup=0 repeat=1 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.011400 median_ms=0.011400 mean_ms=0.011400 effective_bandwidth_gbs=1.080000 max_abs_error=0

三项记录均为 precheck=OK postcheck=OK correct=OK，max_abs_error=0。小输入下程序会收缩 grid，三份输出的实际 grid=1，不是请求上限 256。日志中的单次 event 数值不作性能比较；边界通过后，再回到相同主输入与实际 grid=256 的对照。

运行与校验

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
mkdir -p "$RUN/configs/hip-v3-g256" && \
test ! -e "$RUN/configs/hip-v3-g256/benchmark-p1.samples.csv" && \
"$RUN/build/vector_add_hip" \
  --version v3 --block 256 --grid 256 \
  --size 16777216 --warmup 10 --repeat 50 --seed 20260920 \
  --samples "$RUN/configs/hip-v3-g256/benchmark-p1.samples.csv" \
  --config-id hip-v3-g256 --process 1

运行输出：hip-v3-g256，第 1 个独立进程

text

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 multi_processor_count=32 wave_size=32
RESULT operator=vector-add implementation=hip-v3 runtime=hip shape=16777216 dtype=float32 block=256 grid=256 warmup=10 repeat=50 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.340083 median_ms=0.342002 mean_ms=0.342412 effective_bandwidth_gbs=588.670796 max_abs_error=0

先核对 grid=256、block=256 与标量 v2 一致，前后检查通过；本次单进程中位数为 0.342002 ms。四元素分组同时改变循环工作量、地址表达和编译选择，不能仅凭源码中的类型推断访存事务减少。

相同 grid 和 block 下标量循环与 float4 循环的时间比较

图 8.6　固定 grid=256、block=256，比较 v2 与 v3。相同执行配置使这轮改动更容易解释。

标量 v2 为 0.343640 ms，四元素分组 v3 为 0.342121 ms。两者的三进程范围重叠，与 v0 也有重叠，尚无稳定收益足以替换简单基线。

因此本轮停止这项改动，保留 v0。只有还要回答“编译器是否生成了预期的宽指令”时，才继续看编译结果；即使发现指令变宽，也不能据此宣布速度会提高。8.5.3 的 ATT 选读只演示 v0 的源码与指令对应，不作为这里 v3 的解释证据。

8.4.5 HIP 路线的实验决策

这条路线保留了三种可复用的方法：用相同工作量隔离地址排列，用同一循环 kernel 扫描 grid，用相同 grid/block 对照元素分组。每个对照都回答一个问题；没有稳定提速也能帮助决定停止哪项尝试。

当前主输入选择 v0。换成别的实际长度或布局时，应从基线重新比较，不能把某个 grid 或向量类型当成通用答案。

8.4.6 Triton 基线：一个 program 处理一段元素

HIP 的核函数描述一个线程；Triton 的核函数描述一块数据。我们先让一个 program 负责 256 个连续元素，得到本路线的 triton-t0。它是可以直接修改、复测的基线，不需要先复制 HIP 的线程循环和 float4 写法。

python

@triton.jit
def vector_add_kernel(
    input_a_ptr, input_b_ptr, output_ptr, size,
    BLOCK_SIZE: tl.constexpr,
):
    program_id = tl.program_id(axis=0)
    offsets = program_id * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    valid = offsets < size
    input_a = tl.load(input_a_ptr + offsets, mask=valid, other=0.0)
    input_b = tl.load(input_b_ptr + offsets, mask=valid, other=0.0)
    tl.store(output_ptr + offsets, input_a + input_b, mask=valid)

tl.arange 生成一组下标，program_id * BLOCK_SIZE 将它们移到当前片段。tl.load 读入两组数，加法逐位置进行，tl.store 写回同一组位置。BLOCK_SIZE 表示逻辑元素数，不是 HIP 的线程数 blockDim.x；元素怎样分配给底层执行线程与寄存器，由编译器安排。

先用小输入检查最后一段：

8.4 Triton · Tile · N=13, BLOCK_SIZE=8生成 offsets，再用 mask 关掉越界位置

左右滑动画布，查看完整过程

P0 · 读取 A：offsets 全部有效program 0 生成 offsets=[0…7]，mask 全为 true：从 A 读取 8 个位置。

算法过程示意 · 数字定位到该步完成状态，可拖动进度或播放；动画速度不代表 GPU 耗时。

图 8.7　N=13、教学 tile 大小为 8。第二个 program 生成下标 8–15，mask 只允许读写 8–12。

在图 8.7 中，load 的 other=0.0 为关闭的位置提供占位值；store 仍必须独立传入 mask，才能保护输出边界。

主机端负责把整个数组划分成 program 网格并启动它：

python

size = output.numel()
grid = (triton.cdiv(size, implementation.block_size),)
vector_add_kernel[grid](
    input_a, input_b, output, size,
    BLOCK_SIZE=implementation.block_size,
    num_warps=4,
)

triton.cdiv 向上取整。完整程序还包含输出分配、CPU 参考计算、正确性检查与 GPU event 计时。先读 main() 中的准备与校验，再读 benchmark() 如何反复调用 launch()。ROCm 版 PyTorch 沿用 torch.cuda.Event 和 device="cuda" 这些兼容接口名。

完整源码：code/part2-kernels/chapter8/vector_add_triton.py

python

"""Chapter 8 Element-Wise Vector Add, Triton route.

The kernel is intentionally small. The host program still performs a CPU
reference check before timing, warmup, GPU-event timing, and a postcheck. Book
performance claims are valid only after this script runs on Radeon RX 9070 XT,
ROCm 10.0, and native Ubuntu 24.04.
"""

from __future__ import annotations

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

import numpy as np
import torch
import triton
import triton.language as tl


T0_BLOCK_SIZE = 256
NUM_WARPS = 4


@triton.jit
def vector_add_kernel(
    input_a_ptr,
    input_b_ptr,
    output_ptr,
    size,
    BLOCK_SIZE: tl.constexpr,
):
    """Add one tile of two FP32 vectors per Triton program."""
    program_id = tl.program_id(axis=0)
    offsets = program_id * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    valid = offsets < size

    input_a = tl.load(input_a_ptr + offsets, mask=valid, other=0.0)
    input_b = tl.load(input_b_ptr + offsets, mask=valid, other=0.0)
    tl.store(output_ptr + offsets, input_a + input_b, mask=valid)


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


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


@dataclass(frozen=True)
class Timing:
    minimum_ms: float
    median_ms: float
    mean_ms: float
    samples_ms: list[float]


def parse_args() -> argparse.Namespace:
    parser = argparse.ArgumentParser()
    parser.add_argument("--version", choices=("all", "t0", "t1"), default="all")
    parser.add_argument("--size", type=int, default=16 * 1024 * 1024)
    parser.add_argument("--block", type=int, default=1024)
    parser.add_argument("--warmup", type=int, default=10)
    parser.add_argument("--repeat", type=int, default=50)
    parser.add_argument("--seed", type=int, default=20260716)
    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.size <= 0:
        parser.error("--size must be positive")
    if args.block <= 0 or args.block & (args.block - 1):
        parser.error("--block must be a positive power of two")
    if args.warmup < 0 or args.repeat <= 0:
        parser.error("--warmup must be non-negative and --repeat positive")
    if args.process <= 0 or (args.config_id and not re.fullmatch(r"[A-Za-z0-9_.-]+", args.config_id)):
        parser.error("--process must be positive; --config-id uses letters, digits, '-', '_' or '.'")
    return args


def selected_implementations(version: str, t1_block: int) -> list[Implementation]:
    t0 = Implementation("triton-t0", T0_BLOCK_SIZE)
    t1 = Implementation("triton-t1", t1_block)
    if version == "all":
        return [t0, t1]
    return [t0 if version == "t0" else t1]


def make_host_inputs(size: int, seed: int) -> tuple[torch.Tensor, torch.Tensor]:
    """Generate the same deterministic FP32 values as the HIP program."""
    indices = np.arange(size, dtype=np.uint64)
    seed_u64 = np.uint64(seed & ((1 << 64) - 1))
    with np.errstate(over="ignore"):
        input_a_integer = ((indices * np.uint64(17) + seed_u64) % 2048).astype(
            np.int64
        ) - 1024
        input_b_integer = (
            (indices * np.uint64(29) + seed_u64 * np.uint64(3)) % 4096
        ).astype(np.int64) - 2048
    input_a = torch.from_numpy((input_a_integer / 1024.0).astype(np.float32))
    input_b = torch.from_numpy((input_b_integer / 2048.0).astype(np.float32))
    return input_a, input_b


def launch(
    implementation: Implementation,
    input_a: torch.Tensor,
    input_b: torch.Tensor,
    output: torch.Tensor,
) -> None:
    size = output.numel()
    grid = (triton.cdiv(size, implementation.block_size),)
    vector_add_kernel[grid](
        input_a,
        input_b,
        output,
        size,
        BLOCK_SIZE=implementation.block_size,
        num_warps=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())
    finite = bool(torch.isfinite(actual).all().item())
    return Validation(
        correct=finite and max_absolute_error <= 1.0e-6,
        max_absolute_error=max_absolute_error,
    )


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

    starts = [torch.cuda.Event(enable_timing=True) for _ in range(repeat)]
    ends = [torch.cuda.Event(enable_timing=True) for _ in range(repeat)]
    for start, end in zip(starts, ends, strict=True):
        start.record()
        launch(implementation, input_a, input_b, output)
        end.record()
    torch.cuda.synchronize()

    times_ms = [
        start.elapsed_time(end) for start, end in zip(starts, ends, strict=True)
    ]
    return Timing(
        minimum_ms=min(times_ms),
        median_ms=statistics.median(times_ms),
        mean_ms=statistics.fmean(times_ms),
        samples_ms=times_ms,
    )


def write_samples(implementation: Implementation, timing: Timing, args: argparse.Namespace) -> None:
    if args.samples is None:
        return
    fields = ("config_id", "implementation", "runtime", "shape", "dtype", "block", "grid",
              "num_warps", "warmup", "repeat", "seed", "process", "sample", "ms")
    header = not args.samples.exists() or args.samples.stat().st_size == 0
    with args.samples.open("a", newline="") as file:
        writer = csv.writer(file)
        if header:
            writer.writerow(fields)
        for index, milliseconds in enumerate(timing.samples_ms):
            writer.writerow((args.config_id or implementation.name, implementation.name, "triton",
                             args.size, "float32", implementation.block_size,
                             triton.cdiv(args.size, implementation.block_size), NUM_WARPS,
                             args.warmup, args.repeat, args.seed, args.process, index,
                             format(milliseconds, ".12g")))


def print_failed_result(
    implementation: Implementation,
    validation: Validation,
    args: argparse.Namespace,
) -> None:
    grid = triton.cdiv(args.size, implementation.block_size)
    print(
        "RESULT"
        " operator=vector-add"
        f" implementation={implementation.name}"
        " runtime=triton"
        f" shape={args.size}"
        " dtype=float32"
        f" block={implementation.block_size}"
        f" grid={grid}"
        f" num_warps={NUM_WARPS}"
        f" warmup={args.warmup}"
        f" repeat={args.repeat}"
        f" seed={args.seed}"
        " timed=0 correct=FAIL precheck=FAIL postcheck=NA"
        " min_ms=NA median_ms=NA mean_ms=NA"
        " effective_bandwidth_gbs=NA"
        f" max_abs_error={validation.max_absolute_error:.9g}"
    )


def print_result(
    implementation: Implementation,
    timing: Timing,
    validation: Validation,
    args: argparse.Namespace,
) -> None:
    grid = triton.cdiv(args.size, implementation.block_size)
    logical_bytes = (
        3.0 * args.size * torch.tensor([], dtype=torch.float32).element_size()
    )
    effective_bandwidth_gbs = logical_bytes / (timing.median_ms * 1.0e-3) / 1.0e9
    status = "OK" if validation.correct else "FAIL"
    print(
        "RESULT"
        " operator=vector-add"
        f" implementation={implementation.name}"
        " runtime=triton"
        f" shape={args.size}"
        " dtype=float32"
        f" block={implementation.block_size}"
        f" grid={grid}"
        f" num_warps={NUM_WARPS}"
        f" warmup={args.warmup}"
        f" repeat={args.repeat}"
        f" seed={args.seed}"
        f" timed=1 correct={status} precheck=OK postcheck={status}"
        f" min_ms={timing.minimum_ms:.6f}"
        f" median_ms={timing.median_ms:.6f}"
        f" mean_ms={timing.mean_ms:.6f}"
        f" effective_bandwidth_gbs={effective_bandwidth_gbs:.6f}"
        f" max_abs_error={validation.max_absolute_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_a, host_b = make_host_inputs(args.size, args.seed)
    reference = host_a + host_b
    device_a = host_a.to(device="cuda")
    device_b = host_b.to(device="cuda")
    output = torch.empty_like(device_a)

    all_correct = True
    for implementation in selected_implementations(args.version, args.block):
        output.fill_(float("nan"))
        launch(implementation, device_a, device_b, output)
        torch.cuda.synchronize()
        precheck = validate(output, reference)
        if not precheck.correct:
            print_failed_result(implementation, precheck, args)
            all_correct = False
            continue

        timing = benchmark(
            implementation,
            device_a,
            device_b,
            output,
            warmup=args.warmup,
            repeat=args.repeat,
        )
        postcheck = validate(output, reference)
        if postcheck.correct:
            write_samples(implementation, timing, args)
        print_result(implementation, timing, postcheck, args)
        all_correct = all_correct and postcheck.correct

    return 0 if all_correct else 1


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

运行与校验

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
mkdir -p "$RUN/configs/triton-t0-b256" && \
test ! -e "$RUN/configs/triton-t0-b256/benchmark-p1.samples.csv" && \
python chapter8/vector_add_triton.py \
  --version t0 --block 256 \
  --size 16777216 --warmup 10 --repeat 50 --seed 20260920 \
  --samples "$RUN/configs/triton-t0-b256/benchmark-p1.samples.csv" \
  --config-id triton-t0-b256 --process 1

运行输出：triton-t0-b256，第 1 个独立进程

text

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=vector-add implementation=triton-t0 runtime=triton shape=16777216 dtype=float32 block=256 grid=65536 num_warps=4 warmup=10 repeat=50 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.348681 median_ms=0.349521 mean_ms=0.350191 effective_bandwidth_gbs=576.007123 max_abs_error=0

第一次调用在预检阶段完成 JIT 编译，不进入后面的稳态样本。先确认 correct=OK，再看 block=256、grid=65536、num_warps=4。这里输出字段 block 表示逻辑 tile 的元素数，不是实际线程数。

本次单进程中位数为 0.349521 ms；三进程的中位数为 0.334061 ms，范围 0.333781–0.349521 ms。这组结果是本路线后续比较的起点。

基线分析：256 个元素对应多少底层线程？ 源码中的 tile 是逻辑数据量。下面示范本路线的基本 trace，核对编译后工作组与 program 网格，不用它替代上面的 benchmark。

仍在本篇目录运行；预热 5 次、重复 10 次，另有 1 次预检。inspect_trace.py 匹配目标函数，按时间排序，跳过 6 条预检与预热，只读取后续 10 条。其完整源码在 8.5.1，只依赖标准库。

256 个元素：资源参照。

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
mkdir "$RUN/configs/triton-t0-b256/profile" && \
rocprofv3 --kernel-trace --hip-trace --stats \
  --output-format csv pftrace \
  --output-directory "$RUN/configs/triton-t0-b256/profile" \
  --output-file triton-t0-b256 -- \
  python chapter8/vector_add_triton.py \
  --version t0 --block 256 \
  --size 16777216 --warmup 5 --repeat 10 --seed 20260920 \
  --config-id triton-t0-b256 --process 1 && \
python chapter8/inspect_trace.py \
  "$RUN/configs/triton-t0-b256/profile/triton-t0-b256_kernel_trace.csv" \
  --kernel vector_add_kernel --skip 6 --take 10

性能分析原始输出：triton-t0-b256（终端、生成记录与 CSV）

程序 stdout

text

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=vector-add implementation=triton-t0 runtime=triton shape=16777216 dtype=float32 block=256 grid=65536 num_warps=4 warmup=5 repeat=10 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.335801 median_ms=0.337841 mean_ms=1.341814 effective_bandwidth_gbs=595.920260 max_abs_error=0

Profiler stderr：文件生成记录

text

W0920 08:17:21.035465 1401930 simple_timer.cpp:55] [rocprofv3] tool initialization ::     0.001323 sec
W0920 08:17:21.035661 1401930 simple_timer.cpp:55] [rocprofv3] 'python chapter8/vector_add_triton.py --version t0 --size 16777216 --block 256 --warmup 5 --repeat 10 --seed 20260920 --config-id triton-t0-b256 --process 1' ::     0.000000 sec
W0920 08:17:21.229178 1401930 tool.cpp:3348] HIP (compiler) version 7.15.0 initialized (instance=0)
W0920 08:17:22.079010 1401930 tool.cpp:3348] HIP (runtime) version 7.15.0 initialized (instance=0)
W0920 08:17:22.085151 1401930 tool.cpp:3348] HSA version 1.21.0 initialized (instance=0)
W0920 08:17:22.456215 1401930 queue_controller.cpp:170] [queue-intercept] device-memory ring-buffer requested but profiling requires system-memory InterceptQueue; falling back to system-memory ring for agent 417639552 (priority and CU-mask from the descriptor are not preserved)
W0920 08:17:22.559938 1401990 simple_timer.cpp:55] [rocprofv3] tool initialization ::     0.001680 sec
W0920 08:17:22.560126 1401990 simple_timer.cpp:55] [rocprofv3] '/bin/sh /sbin/ldconfig -p' ::     0.000000 sec
E0920 08:17:23.626694 1401930 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/triton-t0-b256/profile/triton-t0-b256_kernel_stats.csv
E0920 08:17:23.626888 1401930 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/triton-t0-b256/profile/triton-t0-b256_kernel_trace.csv
E0920 08:17:23.648852 1401930 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/triton-t0-b256/profile/triton-t0-b256_hip_api_stats.csv
E0920 08:17:23.649059 1401930 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/triton-t0-b256/profile/triton-t0-b256_hip_api_trace.csv
E0920 08:17:23.741952 1401930 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/triton-t0-b256/profile/triton-t0-b256_agent_info.csv
E0920 08:17:23.742093 1401930 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/triton-t0-b256/profile/triton-t0-b256_domain_stats.csv
[159.793]       perfetto.cc:47314 Configured tracing session 1, #sources:1, duration:0 ms, #buffers:1, total buffer size:1048576 KB, total sessions:1, uid:0 session name: ""
E0920 08:17:24.392122 1401930 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/triton-t0-b256/profile/triton-t0-b256_results.pftrace
[160.479]       perfetto.cc:48889 Tracing session 1 ended, total sessions:0
W0920 08:17:24.439865 1401930 simple_timer.cpp:55] [rocprofv3] output generation ::     0.834389 sec
W0920 08:17:24.440022 1401930 simple_timer.cpp:55] [rocprofv3] tool finalization ::     0.834609 sec

原始 kernel trace CSV

text

"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",2,0,1401930,1,56,"void at::native::vectorized_elementwise_kernel<4, at::native::FillFunctor<float>, std::array<char*, 1ul> >(int, at::native::FillFunctor<float>, std::array<char*, 1ul>)",98200,908158564559237,908158564620117,0,0,8,0,128,256,1,1,4194304,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401930,2,187,"vector_add_kernel",98230,908159283436218,908159285827907,0,0,8,0,128,128,1,1,8388608,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401930,3,187,"vector_add_kernel",98247,908159343374101,908159343835862,0,0,8,0,128,128,1,1,8388608,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401930,4,187,"vector_add_kernel",98252,908159343841062,908159344307624,0,0,8,0,128,128,1,1,8388608,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401930,5,187,"vector_add_kernel",98257,908159344313184,908159344967987,0,0,8,0,128,128,1,1,8388608,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401930,6,187,"vector_add_kernel",98262,908159344982227,908159345345868,0,0,8,0,128,128,1,1,8388608,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401930,7,187,"vector_add_kernel",98267,908159345351268,908159345714590,0,0,8,0,128,128,1,1,8388608,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401930,8,187,"vector_add_kernel",98286,908159345850020,908159346359142,0,0,8,0,128,128,1,1,8388608,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401930,11,187,"vector_add_kernel",98361,908159347060465,908159347392226,0,0,8,0,128,128,1,1,8388608,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401930,10,187,"vector_add_kernel",98336,908159346716184,908159347044265,0,0,8,0,128,128,1,1,8388608,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401930,9,187,"vector_add_kernel",98311,908159346374782,908159346700223,0,0,8,0,128,128,1,1,8388608,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401930,12,187,"vector_add_kernel",98386,908159357237336,908159357564737,0,0,8,0,128,128,1,1,8388608,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401930,13,187,"vector_add_kernel",98411,908159357580257,908159357905858,0,0,8,0,128,128,1,1,8388608,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401930,14,187,"vector_add_kernel",98436,908159357921498,908159358248979,0,0,8,0,128,128,1,1,8388608,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401930,15,187,"vector_add_kernel",98461,908159358264619,908159358590020,0,0,8,0,128,128,1,1,8388608,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401930,16,187,"vector_add_kernel",98486,908159358605500,908159358932662,0,0,8,0,128,128,1,1,8388608,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1401930,17,187,"vector_add_kernel",98511,908159358948262,908159359275343,0,0,8,0,128,128,1,1,8388608,1,1

筛选后的分析结果：triton-t0-b256

text

kernel_name="vector_add_kernel"
csv_rows=17 target_rows=16 skipped_rows=6 kept_rows=10
sort=Start_Timestamp_ascending
first_start_ns=908159345850020 first_end_ns=908159346359142 first_duration_ns=509122
kernel_min_us=325.401000 kernel_median_us=327.281500 kernel_max_us=509.122000
grid_work_items_xyz=8388608x1x1 workgroup_size_xyz=128x1x1 workgroups_xyz=65536x1x1
workgroups_total=65536
vgpr_count=8 sgpr_count=128 lds_bytes=0 scratch_bytes=0

先核对 target_rows=16 与 kept_rows=10。Workgroup X 为 128；Grid X 为 8,388,608，两者相除是 65,536 个 workgroup，与本次 program 数一致。可见，256 个逻辑元素并不等于 256 个底层线程。记录 vgpr_count=8，后续如有资源问题可复用这份基线，不必重新采集。

时间字段是设备 kernel 的起止时间，单位为 μs；分析窗口中位数为 327.2815 μs。它与前面的 event 是不同运行、不同范围，不能相减来推算 CPU 开销。当前已核对到执行映射，可以继续检验更大的逻辑 tile。

现在可以直接提出分工假设：一个 program 多处理一些元素，减少 program 数量，是否有收益？先固定 num_warps=4 比较 tile，保留原有连续访问与 mask，不需要照搬 HIP 的 grid-stride 或 float4。

8.4.7 Tile 扫描：固定执行参数，只改数据粒度

先比较四个候选，观察每项实际改变了什么：

每 program 的元素数 `BLOCK_SIZE`	program 数，`N=16,777,216`	`num_warps`
256	65,536	4
512	32,768	4
1,024	16,384	4
2,048	8,192	4

源码中的 t0 固定使用 256 个元素；扫描参数时选择 t1，再用 --block 指定逻辑 tile。三项都调用完整源码中的 vector_add_kernel()；launch() 依据 tile 计算网格，改的是传入编译器的配置。

四项使用同一个 kernel，地址仍连续，尾部仍由 mask 保护。num_warps 控制编译一个 program 时采用的执行组数；它与“这一份工作有多少元素”是不同的参数。本轮固定为 4，以便先理解 tile 大小这一项选择。

更大的 tile 同时带来两种变化：独立 program 更少，每个 program 需要处理和保留的数据更多。 因而不能只计算 program 数减少了几倍，就推出性能也会按比例提高。

改变 tile 后，先检查各候选的「少一个、恰好填满、多一个」。这样可以验证 mask 与 program 网格配合时不会漏掉最后一个元素；这些短运行不用于测速。

bash

python chapter8/vector_add_triton.py --version t1 --size 511 --block 512 --warmup 0 --repeat 1 --seed 20260920
python chapter8/vector_add_triton.py --version t1 --size 512 --block 512 --warmup 0 --repeat 1 --seed 20260920
python chapter8/vector_add_triton.py --version t1 --size 513 --block 512 --warmup 0 --repeat 1 --seed 20260920
python chapter8/vector_add_triton.py --version t1 --size 1023 --block 1024 --warmup 0 --repeat 1 --seed 20260920
python chapter8/vector_add_triton.py --version t1 --size 1024 --block 1024 --warmup 0 --repeat 1 --seed 20260920
python chapter8/vector_add_triton.py --version t1 --size 1025 --block 1024 --warmup 0 --repeat 1 --seed 20260920
python chapter8/vector_add_triton.py --version t1 --size 2047 --block 2048 --warmup 0 --repeat 1 --seed 20260920
python chapter8/vector_add_triton.py --version t1 --size 2048 --block 2048 --warmup 0 --repeat 1 --seed 20260920
python chapter8/vector_add_triton.py --version t1 --size 2049 --block 2048 --warmup 0 --repeat 1 --seed 20260920

边界检查原始输出：三个 tile 的边界 @ RX 9070 XT / ROCm 10.0

boundary-tile512-n511.log

text

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=vector-add implementation=triton-t1 runtime=triton shape=511 dtype=float32 block=512 grid=1 num_warps=4 warmup=0 repeat=1 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.090331 median_ms=0.090331 mean_ms=0.090331 effective_bandwidth_gbs=0.067884 max_abs_error=0

boundary-tile512-n512.log

text

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=vector-add implementation=triton-t1 runtime=triton shape=512 dtype=float32 block=512 grid=1 num_warps=4 warmup=0 repeat=1 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.076541 median_ms=0.076541 mean_ms=0.076541 effective_bandwidth_gbs=0.080271 max_abs_error=0

boundary-tile512-n513.log

text

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=vector-add implementation=triton-t1 runtime=triton shape=513 dtype=float32 block=512 grid=2 num_warps=4 warmup=0 repeat=1 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.122852 median_ms=0.122852 mean_ms=0.122852 effective_bandwidth_gbs=0.050109 max_abs_error=0

boundary-tile1024-n1023.log

text

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=vector-add implementation=triton-t1 runtime=triton shape=1023 dtype=float32 block=1024 grid=1 num_warps=4 warmup=0 repeat=1 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.080341 median_ms=0.080341 mean_ms=0.080341 effective_bandwidth_gbs=0.152799 max_abs_error=0

boundary-tile1024-n1024.log

text

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=vector-add implementation=triton-t1 runtime=triton shape=1024 dtype=float32 block=1024 grid=1 num_warps=4 warmup=0 repeat=1 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.083081 median_ms=0.083081 mean_ms=0.083081 effective_bandwidth_gbs=0.147904 max_abs_error=0

boundary-tile1024-n1025.log

text

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=vector-add implementation=triton-t1 runtime=triton shape=1025 dtype=float32 block=1024 grid=2 num_warps=4 warmup=0 repeat=1 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.088171 median_ms=0.088171 mean_ms=0.088171 effective_bandwidth_gbs=0.139502 max_abs_error=0

boundary-tile2048-n2047.log

text

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=vector-add implementation=triton-t1 runtime=triton shape=2047 dtype=float32 block=2048 grid=1 num_warps=4 warmup=0 repeat=1 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.077201 median_ms=0.077201 mean_ms=0.077201 effective_bandwidth_gbs=0.318182 max_abs_error=0

boundary-tile2048-n2048.log

text

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=vector-add implementation=triton-t1 runtime=triton shape=2048 dtype=float32 block=2048 grid=1 num_warps=4 warmup=0 repeat=1 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.074481 median_ms=0.074481 mean_ms=0.074481 effective_bandwidth_gbs=0.329963 max_abs_error=0

boundary-tile2048-n2049.log

text

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=vector-add implementation=triton-t1 runtime=triton shape=2049 dtype=float32 block=2048 grid=2 num_warps=4 warmup=0 repeat=1 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.089301 median_ms=0.089301 mean_ms=0.089301 effective_bandwidth_gbs=0.275338 max_abs_error=0

九项均通过前后检查，最大绝对误差为零；每组三个长度的 grid 分别为 1、1、2。不要用这些无预热、单次执行的 event 时间给候选排序。下面恢复相同主输入与正式预热、重复次数，比较 tile 的性能。

每个 program 处理 512 个元素。

运行与校验

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
mkdir -p "$RUN/configs/triton-t1-b512" && \
test ! -e "$RUN/configs/triton-t1-b512/benchmark-p1.samples.csv" && \
python chapter8/vector_add_triton.py \
  --version t1 --block 512 \
  --size 16777216 --warmup 10 --repeat 50 --seed 20260920 \
  --samples "$RUN/configs/triton-t1-b512/benchmark-p1.samples.csv" \
  --config-id triton-t1-b512 --process 1

运行输出：triton-t1-b512，第 1 个独立进程

text

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=vector-add implementation=triton-t1 runtime=triton shape=16777216 dtype=float32 block=512 grid=32768 num_warps=4 warmup=10 repeat=50 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.340041 median_ms=0.340481 mean_ms=0.340498 effective_bandwidth_gbs=591.300496 max_abs_error=0

核对 block=512、grid=32768，以及三个检查字段为 OK。本次单进程中位数为 0.340481 ms。

每个 program 处理 1,024 个元素。

运行与校验

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
mkdir -p "$RUN/configs/triton-t1-b1024" && \
test ! -e "$RUN/configs/triton-t1-b1024/benchmark-p1.samples.csv" && \
python chapter8/vector_add_triton.py \
  --version t1 --block 1024 \
  --size 16777216 --warmup 10 --repeat 50 --seed 20260920 \
  --samples "$RUN/configs/triton-t1-b1024/benchmark-p1.samples.csv" \
  --config-id triton-t1-b1024 --process 1

运行输出：triton-t1-b1024，第 1 个独立进程

text

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=vector-add implementation=triton-t1 runtime=triton shape=16777216 dtype=float32 block=1024 grid=16384 num_warps=4 warmup=10 repeat=50 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.337521 median_ms=0.338220 mean_ms=0.338648 effective_bandwidth_gbs=595.252496 max_abs_error=0

核对 block=1024、grid=16384。本次单进程中位数为 0.338220 ms；后面仍与最初的 tile256 一起比较，不只挑一个较慢候选作分母。

每个 program 处理 2,048 个元素。

运行与校验

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
mkdir -p "$RUN/configs/triton-t1-b2048" && \
test ! -e "$RUN/configs/triton-t1-b2048/benchmark-p1.samples.csv" && \
python chapter8/vector_add_triton.py \
  --version t1 --block 2048 \
  --size 16777216 --warmup 10 --repeat 50 --seed 20260920 \
  --samples "$RUN/configs/triton-t1-b2048/benchmark-p1.samples.csv" \
  --config-id triton-t1-b2048 --process 1

运行输出：triton-t1-b2048，第 1 个独立进程

text

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=vector-add implementation=triton-t1 runtime=triton shape=16777216 dtype=float32 block=2048 grid=8192 num_warps=4 warmup=10 repeat=50 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.340561 median_ms=0.341241 mean_ms=0.341815 effective_bandwidth_gbs=589.983592 max_abs_error=0

核对 block=2048、grid=8192。本次单进程中位数为 0.341241 ms。program 更少没有在这一次运行里自动换来更短的时间，判断仍看三进程结果。

同一 Triton kernel 在固定 num_warps 下的四种 tile 大小与时间

图 8.8　固定 num_warps=4 的 tile 扫描；每个候选均先通过正确性，再参加独立进程计时。

256、512、1,024、2,048 四个候选分别为 0.334061、0.335441、0.337701、0.341241 ms。三个更大 tile 的进程范围都与 256 基线有重叠，没有稳定收益，因此本轮保留 BLOCK_SIZE=256、num_warps=4。这个性能决定到这里已经成立。

代表配置分析：program 减少后，每份工作的资源分配是否也变化？ 如果要进一步理解 tile 的代价，可以比较扫描的两端 256 与 2048。选择它们是为了检查分工变化，不是为了把“没有收益”解释成已经证明的寄存器瓶颈；中间两个配置无需再重复同一操作。

复用 8.4.6 中已取得的 tile256 基线，只对 tile2048 增加一次分析。下面仍在本篇目录运行，采用同样的预检、预热与筛选窗口；两份 trace 都独立于 benchmark。

2,048 个元素：与参照保持相同 num_warps。

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
mkdir "$RUN/configs/triton-t1-b2048/profile" && \
rocprofv3 --kernel-trace --hip-trace --stats \
  --output-format csv pftrace \
  --output-directory "$RUN/configs/triton-t1-b2048/profile" \
  --output-file triton-t1-b2048 -- \
  python chapter8/vector_add_triton.py \
  --version t1 --block 2048 \
  --size 16777216 --warmup 5 --repeat 10 --seed 20260920 \
  --config-id triton-t1-b2048 --process 1 && \
python chapter8/inspect_trace.py \
  "$RUN/configs/triton-t1-b2048/profile/triton-t1-b2048_kernel_trace.csv" \
  --kernel vector_add_kernel --skip 6 --take 10

性能分析原始输出：triton-t1-b2048（终端、生成记录与 CSV）

程序 stdout

text

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=vector-add implementation=triton-t1 runtime=triton shape=16777216 dtype=float32 block=2048 grid=8192 num_warps=4 warmup=5 repeat=10 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.344761 median_ms=0.345321 mean_ms=1.437385 effective_bandwidth_gbs=583.012884 max_abs_error=0

Profiler stderr：文件生成记录

text

W0920 08:17:32.010878 1402203 simple_timer.cpp:55] [rocprofv3] tool initialization ::     0.001075 sec
W0920 08:17:32.011063 1402203 simple_timer.cpp:55] [rocprofv3] 'python chapter8/vector_add_triton.py --version t1 --size 16777216 --block 2048 --warmup 5 --repeat 10 --seed 20260920 --config-id triton-t1-b2048 --process 1' ::     0.000000 sec
W0920 08:17:32.206664 1402203 tool.cpp:3348] HIP (compiler) version 7.15.0 initialized (instance=0)
W0920 08:17:33.067973 1402203 tool.cpp:3348] HIP (runtime) version 7.15.0 initialized (instance=0)
W0920 08:17:33.074098 1402203 tool.cpp:3348] HSA version 1.21.0 initialized (instance=0)
W0920 08:17:33.423933 1402203 queue_controller.cpp:170] [queue-intercept] device-memory ring-buffer requested but profiling requires system-memory InterceptQueue; falling back to system-memory ring for agent 595855312 (priority and CU-mask from the descriptor are not preserved)
W0920 08:17:33.526777 1402263 simple_timer.cpp:55] [rocprofv3] tool initialization ::     0.001023 sec
W0920 08:17:33.526941 1402263 simple_timer.cpp:55] [rocprofv3] '/bin/sh /sbin/ldconfig -p' ::     0.000000 sec
E0920 08:17:34.602277 1402203 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/triton-t1-b2048/profile/triton-t1-b2048_kernel_stats.csv
E0920 08:17:34.602472 1402203 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/triton-t1-b2048/profile/triton-t1-b2048_kernel_trace.csv
E0920 08:17:34.623736 1402203 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/triton-t1-b2048/profile/triton-t1-b2048_hip_api_stats.csv
E0920 08:17:34.623948 1402203 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/triton-t1-b2048/profile/triton-t1-b2048_hip_api_trace.csv
E0920 08:17:34.715478 1402203 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/triton-t1-b2048/profile/triton-t1-b2048_agent_info.csv
E0920 08:17:34.715613 1402203 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/triton-t1-b2048/profile/triton-t1-b2048_domain_stats.csv
[170.766]       perfetto.cc:47314 Configured tracing session 1, #sources:1, duration:0 ms, #buffers:1, total buffer size:1048576 KB, total sessions:1, uid:0 session name: ""
E0920 08:17:35.384481 1402203 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/configs/triton-t1-b2048/profile/triton-t1-b2048_results.pftrace
[171.472]       perfetto.cc:48889 Tracing session 1 ended, total sessions:0
W0920 08:17:35.432395 1402203 simple_timer.cpp:55] [rocprofv3] output generation ::     0.849732 sec
W0920 08:17:35.432548 1402203 simple_timer.cpp:55] [rocprofv3] tool finalization ::     0.850157 sec

原始 kernel trace CSV

text

"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",2,0,1402203,1,112,"void at::native::vectorized_elementwise_kernel<4, at::native::FillFunctor<float>, std::array<char*, 1ul> >(int, at::native::FillFunctor<float>, std::array<char*, 1ul>)",98200,908169531682554,908169531743995,0,0,8,0,128,256,1,1,4194304,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1402203,2,187,"vector_add_kernel",98230,908170251288615,908170254234346,0,0,40,0,128,128,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1402203,3,187,"vector_add_kernel",98247,908170311206504,908170311674385,0,0,40,0,128,128,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1402203,4,187,"vector_add_kernel",98252,908170311679545,908170312152587,0,0,40,0,128,128,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1402203,5,187,"vector_add_kernel",98257,908170312158507,908170312631069,0,0,40,0,128,128,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1402203,6,187,"vector_add_kernel",98262,908170312636989,908170313292872,0,0,40,0,128,128,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1402203,7,187,"vector_add_kernel",98267,908170313307512,908170313678593,0,0,40,0,128,128,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1402203,8,187,"vector_add_kernel",98286,908170313989815,908170314521377,0,0,40,0,128,128,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1402203,9,187,"vector_add_kernel",98311,908170325154133,908170325489014,0,0,40,0,128,128,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1402203,10,187,"vector_add_kernel",98336,908170325505294,908170325839175,0,0,40,0,128,128,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1402203,11,187,"vector_add_kernel",98361,908170325855375,908170326189576,0,0,40,0,128,128,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1402203,12,187,"vector_add_kernel",98386,908170326205816,908170326540418,0,0,40,0,128,128,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1402203,13,187,"vector_add_kernel",98411,908170326556418,908170326890699,0,0,40,0,128,128,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1402203,14,187,"vector_add_kernel",98436,908170326906939,908170327240860,0,0,40,0,128,128,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1402203,15,187,"vector_add_kernel",98461,908170327256700,908170327590622,0,0,40,0,128,128,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1402203,16,187,"vector_add_kernel",98486,908170327606502,908170327940823,0,0,40,0,128,128,1,1,1048576,1,1
"KERNEL_DISPATCH","Agent 1",2,0,1402203,17,187,"vector_add_kernel",98511,908170327956703,908170328291265,0,0,40,0,128,128,1,1,1048576,1,1

筛选后的分析结果：triton-t1-b2048

text

kernel_name="vector_add_kernel"
csv_rows=17 target_rows=16 skipped_rows=6 kept_rows=10
sort=Start_Timestamp_ascending
first_start_ns=908170313989815 first_end_ns=908170314521377 first_duration_ns=531562
kernel_min_us=333.881000 kernel_median_us=334.301000 kernel_max_us=531.562000
grid_work_items_xyz=1048576x1x1 workgroup_size_xyz=128x1x1 workgroups_xyz=8192x1x1
workgroups_total=8192
vgpr_count=40 sgpr_count=128 lds_bytes=0 scratch_bytes=0

阅读时先确认两份结果都为 target_rows=16、kept_rows=10，随后按下列顺序比较：

字段	tile256	tile2048	能回答什么
Workgroup X	128	128	本次实际工作组大小没有随逻辑 tile 变成 8 倍
Grid X	8,388,608	1,048,576	工作项数；除以 Workgroup X，得到 65,536 与 8,192 个 workgroup，与 program 网格对应
VGPR	8	40	当前编译与启动的寄存器资源字段增大
LDS / Scratch（Byte）	0 / 0	0 / 0	本次记录的这两项分配为零

这组证据说明，减少 program 并不是只减少启动数量：每份工作的资源需求也会改变。它没有测出动态 occupancy 或寄存器造成的耗时比例；scratch=0 也不表示资源没有代价。不要将“VGPR 增大”和“没有稳定加速”直接拼成唯一因果关系。

当前仍保留 tile256。若实际工作需要处理另一种长度，可以重新比较；若继续研究资源与并行量的取舍，可固定 tile 再单独改变 num_warps，但那是下一项实验，不能用本轮数据预告结果。

8.4.8 从手工比较到自动调优

Triton 的优化从 program 的数据分工出发：先写清 offsets、mask 与负责的输出，再固定一个执行参数，扫描少量有理由的 tile。编译器负责将逻辑数据块映射到底层线程和寄存器；需要解释这层映射时才查看相应记录。

当候选变多、不同 shape 需要不同配置时，可以让 triton.autotune 搜索。但它只比较给定候选，不替代正确性检查、数据流设计或结果解释。本章保留小范围手工扫描；后面的矩阵乘再研究多个 tile 参数怎样组合。

8.5 按问题选择性能分析

打开 profiler 之前，先写一句它要回答的问题，再决定读什么字段。若报告无法改变下一步的选择，就不必为了流程完整而采集。

8.5.1 从问题到字段

当前问题	优先证据	本章的使用位置
修改后是否正确，是否值得替换基线？	正确性检查与独立进程 benchmark	每轮必做；不需要 profiler 才能作取舍
运行的是预期 kernel 和启动规模吗？	kernel 名称、Grid / Workgroup、目标调用次数	HIP 基线 trace
kernel 在设备上何时开始、何时结束？	起止时间戳与 CPU/GPU 时间线	基线示范；后续有提交节奏疑问时再看
tile 改变后资源分配有变化吗？	代表配置的 Workgroup、VGPR、LDS、Scratch	Triton256 与2048配对分析
物理流量或缓存命中率是否变化？	经过有效性核验的相应硬件计数器	本机本次没有取得足够证据，不能用逻辑字节替代
源码对应哪些指令，局部 wave 在哪里等待？	编译结果或局部指令跟踪	进阶诊断；不能直接分摊整个 kernel 的时间

本章的 inspect_trace.py 先匹配一个确定的 kernel 名称，按起点排序，再去掉预检与预热。它会检查记录数量、必需列和各次启动字段是否一致；不满足预期时直接报错。命令中的 grid 仍要与结果人工对照。

完整分析脚本：code/part2-kernels/chapter8/inspect_trace.py

python

"""Read one rocprofv3 kernel-trace window using only Python's standard library."""
from __future__ import annotations

import argparse
import csv
import json
import math
from pathlib import Path
import re
import statistics


GRID_FIELDS = tuple(f"Grid_Size_{axis}" for axis in "XYZ")
GROUP_FIELDS = tuple(f"Workgroup_Size_{axis}" for axis in "XYZ")
RESOURCE_FIELDS = ("VGPR_Count", "SGPR_Count", "LDS_Block_Size", "Scratch_Size")
INTEGER_FIELDS = ("Start_Timestamp", "End_Timestamp", *GRID_FIELDS,
                  *GROUP_FIELDS, *RESOURCE_FIELDS)
REQUIRED_FIELDS = ("Kernel_Name", *INTEGER_FIELDS)


class TraceError(ValueError):
    """The CSV cannot support the requested comparison."""


def analyze_trace(path: Path, kernel: str, skip: int, take: int) -> dict:
    if not kernel.strip():
        raise TraceError("--kernel must be a non-empty name substring")
    if skip < 0 or take <= 0:
        raise TraceError("--skip must be non-negative and --take must be positive")
    rows = []
    with path.open(encoding="utf-8-sig", newline="") as file:
        reader = csv.DictReader(file, strict=True)
        fields = reader.fieldnames or []
        if len(fields) != len(set(fields)):
            raise TraceError("CSV has duplicate column names")
        missing = sorted(set(REQUIRED_FIELDS) - set(fields))
        if missing:
            raise TraceError("CSV is missing required columns: " + ", ".join(missing))
        for row in reader:
            where = f"CSV line {reader.line_num}"
            if None in row or any(value is None for value in row.values()):
                raise TraceError(f"{where}: row width does not match the header")
            if not row["Kernel_Name"].strip():
                raise TraceError(f"{where}: Kernel_Name is empty")
            for field in INTEGER_FIELDS:
                value = row[field].strip()
                if not re.fullmatch(r"[0-9]+", value):
                    raise TraceError(f"{where}: {field} must be a non-negative integer, got {value!r}")
                row[field] = int(value)
            if row["End_Timestamp"] < row["Start_Timestamp"]:
                raise TraceError(f"{where}: End_Timestamp is earlier than Start_Timestamp")
            for grid_field, group_field in zip(GRID_FIELDS, GROUP_FIELDS):
                grid, group = row[grid_field], row[group_field]
                if grid <= 0 or group <= 0 or grid % group:
                    raise TraceError(f"{where}: {grid_field} must be a positive multiple of {group_field}")
            rows.append(row)

    targets = [row for row in rows if kernel in row["Kernel_Name"]]
    names = sorted({row["Kernel_Name"] for row in targets})
    if len(names) != 1:
        detail = "no kernel names" if not names else ", ".join(repr(name) for name in names)
        raise TraceError(f"--kernel must match exactly one distinct Kernel_Name; matched {detail}")
    expected = skip + take
    if len(targets) != expected:
        raise TraceError(f"expected exactly skip + take = {expected} target rows, found {len(targets)}")
    targets.sort(key=lambda row: row["Start_Timestamp"])
    first = targets[0]
    for row in targets[1:]:
        changed = [field for field in (*GRID_FIELDS, *GROUP_FIELDS, *RESOURCE_FIELDS)
                   if row[field] != first[field]]
        if changed:
            raise TraceError("target rows have different launch/resource fields: " + ", ".join(changed))
    kept = targets[skip:]
    durations = [row["End_Timestamp"] - row["Start_Timestamp"] for row in kept]
    grid = [first[field] for field in GRID_FIELDS]
    group = [first[field] for field in GROUP_FIELDS]
    groups = [g // w for g, w in zip(grid, group)]
    return {
        "csv_rows": len(rows), "kernel_name": names[0], "target_rows": len(targets),
        "skipped_rows": skip, "kept_rows": len(kept),
        "first_kept": {"start_ns": kept[0]["Start_Timestamp"],
                       "end_ns": kept[0]["End_Timestamp"], "duration_ns": durations[0]},
        "kernel_time_us": {"min": min(durations) / 1000,
                           "median": statistics.median(durations) / 1000,
                           "max": max(durations) / 1000},
        "launch": {"grid_work_items_xyz": grid, "workgroup_size_xyz": group,
                   "workgroups_xyz": groups, "workgroups_total": math.prod(groups)},
        "resources": dict(zip(("vgpr_count", "sgpr_count", "lds_bytes", "scratch_bytes"),
                              (first[field] for field in RESOURCE_FIELDS))),
    }


def print_report(report: dict) -> None:
    print("kernel_name=" + json.dumps(report["kernel_name"], ensure_ascii=False))
    print(" ".join(f"{key}={report[key]}" for key in
                   ("csv_rows", "target_rows", "skipped_rows", "kept_rows")))
    print("sort=Start_Timestamp_ascending")
    print(" ".join(f"first_{key}={value}" for key, value in report["first_kept"].items()))
    print(" ".join(f"kernel_{key}_us={value:.6f}" for key, value in report["kernel_time_us"].items()))
    launch = report["launch"]
    print(" ".join(f"{key}=" + "x".join(map(str, launch[key])) for key in
                   ("grid_work_items_xyz", "workgroup_size_xyz", "workgroups_xyz")))
    print(f"workgroups_total={launch['workgroups_total']}")
    print(" ".join(f"{key}={value}" for key, value in report["resources"].items()))


def main(argv: list[str] | None = None) -> int:
    parser = argparse.ArgumentParser(description=__doc__)
    parser.add_argument("path", type=Path, help="rocprofv3 *_kernel_trace.csv")
    parser.add_argument("--kernel", required=True, help="case-sensitive Kernel_Name substring")
    parser.add_argument("--skip", required=True, type=int, help="target rows before the measured window")
    parser.add_argument("--take", required=True, type=int, help="exact number of measured target rows")
    parser.add_argument("--json", action="store_true", help="print the same report as JSON")
    args = parser.parse_args(argv)
    try:
        report = analyze_trace(args.path, args.kernel, args.skip, args.take)
    except (TraceError, OSError, UnicodeError, csv.Error) as error:
        parser.error(str(error))
    if args.json:
        print(json.dumps(report, indent=2, ensure_ascii=False))
    else:
        print_report(report)
    return 0


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

采集命令用 mkdir ... && rocprofv3 ... && python ... 连接。目录已存在或采集失败时停止后续步骤，避免覆盖旧记录或误读旧报告。采集过程中可能改变队列处理方式，因此其 event 时间始终与 benchmark 分开。

文件	阅读用途
`*_kernel_trace.csv`	逐条核对目标名称、时间戳和启动字段
`*_hip_api_trace.csv`	CPU 侧 API 调用，用 `Correlation_Id` 关联提交与设备执行
`*_kernel_stats.csv`	整次进程汇总，含预检与预热，不作正式成绩
`*_results.pftrace`	查看 CPU/GPU 时间线与调用关系

工具的 Opened result file 记录位于 stderr，行首带 E；这些行报告文件生成。是否成功还要看退出状态、文件内容与目标记录，不能只看日志前缀。时间戳单位为 ns，分析脚本为 μs，benchmark 为 ms：1 ms = 1,000 μs = 1,000,000 ns。

8.5.2 时间线与工具适用范围

先让工具回答一个具体问题：要知道哪次 kernel 慢，读时间线；要核对执行了多少工作，尝试有效的硬件计数器；要观察局部 wave 的指令轨迹，再进入 ATT。工具越底层，需要核对的采集条件也越多。

本章环境中的 rocprofv3 来自 ROCprofiler-SDK 1.3.5，工具自报 ROCm 10.0.0。ROCm 发行版、SDK 工具和 HIP 组件各有自己的版本号，不能要求它们都显示 10.0。

bash

rocprofv3 --version

查看本机 rocprofv3 版本输出

ROCm wheel 环境中的工具

text

             version: 1.3.5
        git_revision: 6b0e43f341195e203754e08f850e437ff2fc09f9
        library_arch: x86_64-unknown-linux-gnu
         system_name: Linux
    system_processor: x86_64
      system_version: 6.12.73-95.123.amzn2023.x86_64
         compiler_id: Clang
    compiler_version: 23.0.0
        rocm_version: 10.0.0

输出中的 version 与 rocm_version 对应工具版本和 ROCm 版本；system_version 是工具构建信息，实验机系统以本章的环境记录为准。

工具或功能	在本章 RX 9070 XT 上的情况	适合回答的问题
`rocprofv3` kernel / HIP trace	本章 11 个配置均已采集 CSV 与 Perfetto 时间线	CPU 提交了什么，GPU 何时执行，实际启动规模是多少
`rocprofv3` 硬件计数器（PMC）	wave 数取得可核对结果；本次缓存与流量指标返回零，不能用来解释访存效率	指定 dispatch 的某类事件发生了多少次；先核对指标与采集是否有效
`rocprofv3 --att` 指令跟踪	已取得单次 baseline dispatch 的原始跟踪与解码产物	选定硬件位置上的 wave 执行了哪些指令；适合继续定位局部等待
ROCprofiler Systems	官方 ROCm 10.0 矩阵列出 Radeon 支持；本章未实测此工具	更大的应用中，CPU 调用栈、线程活动与 GPU 工作怎样关联
ROCprofiler Compute（原 Omniperf）	当前官方兼容表未列 RX 9070 XT	在其支持设备上使用硬件指标面板；本章不提供该工具的实操命令
Radeon GPU Profiler（RGP）	官方 Linux Ubuntu 24.04 条目为 Vulkan only，不能直接套到这里的 HIP 程序	需要同时符合 GPU、API 与操作系统要求，不能只看显卡型号

支持范围分别见 SDK tracing 文档、ROCm 10.0 兼容矩阵、Compute 支持表和 RGP Requirements。下面只展开本机已经验证过的路径。

先看时间线。 在 Perfetto 中选择 Open trace file，打开本次配置目录内的 *_results.pftrace。先找到 CPU 的 HIP API 轨道与 GPU kernel 轨道，再搜索目标函数名：CPU 的 launch 调用结束，不代表 GPU kernel 已经完成。展开进程轨道，搜索 vector_add_v0，可以依次定位 16 次目标调用；第一条是预检。点选 GPU 区间，在 Details 中看 Duration，在 Preceding Flows 中找到对应的 hipLaunchKernel；比较时仍应使用 8.4 中经过筛选的多次执行，而非随手选中一条。

同一个 profile 同时导出了 API trace 和 kernel trace，可以借助 Correlation_Id 对应一次提交与设备执行。时间线适合找执行先后和空隙，不能仅凭某段空白就断言 CPU 算力不足；同步、依赖和提交节奏都需要结合程序检查。当前命令没有开启 memory-copy trace，因此不把缺少拷贝轨道解释为没有发生数据复制。

本次 Perfetto 导入出现了一项 track_descriptor_conflicting_reservation 提示，目标 kernel 的 16 条记录仍可检索，数量与 CSV 一致。因此时间线用于观察调用关系，本章的筛选统计继续以原始 CSV 为准；不把查看器的导入提示当作 kernel 性能问题。

8.5.3 进阶诊断：计数器与指令跟踪

前面的配置选择不依赖下面两项诊断。只有遇到对应问题时才展开；不能因为工具更底层，就默认它更适合回答当前问题。

计数器能否作为解释依据？ 当前 wave 计数能与启动规模核对，但缓存与流量计数返回零，不足以解释地址对照的时间差。这里的第一步是确认测到的量有意义。

选读：核对 wave 计数，以及缓存与流量指标的有效性

计数器记录某类硬件事件的次数。先选择一个能用程序规模核对的量，比一开始就读几十个百分比更容易。本例选择 SQ_WAVES_sum，核对 baseline 启动的 wave 数。

下面使用 HIP v0 做独立工具核验。若前面只走了 Triton 路线，也先执行这里的编译命令；它生成的仍是 8.4.1 的完整程序。然后查询指标组合，采集 v0 预检和预热之后的 10 次执行：

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
if [[ ! -x "$RUN/build/vector_add_hip" ]]; then
  hipcc --offload-arch=gfx1201 -O3 -std=c++17 \
    chapter8/vector_add_hip.hip -o "$RUN/build/vector_add_hip"
fi && \
rocprofv3-avail -d 0 pmc-check SQ_WAVES_sum GRBM_GUI_ACTIVE && \
mkdir "$RUN/pmc-waves" && \
rocprofv3 --pmc SQ_WAVES_sum GRBM_GUI_ACTIVE \
  --kernel-include-regex vector_add_v0 --kernel-iteration-range 7-16 \
  --output-format csv --output-directory "$RUN/pmc-waves" \
  --output-file baseline -- \
  "$RUN/build/vector_add_hip" --version v0 --size 16777216 --block 256 \
  --warmup 5 --repeat 10 --seed 20260920

-d 0 选择当前查询到的设备；pmc-check 检查这组指标能否一起采集，最终还要看实际记录。--kernel-iteration-range 7-16 在这里筛选 计数器采集的目标调用；它不会替 8.4 的 kernel trace 删除预检与预热，所以 trace 仍然需要自己的筛选步骤。

查看计数器组合检查、运行输出与完整原始 CSV

组合检查

text

Following input counters can be collected together on GPU:0	SQ_WAVES_sum	GRBM_GUI_ACTIVE

采集进程的输出（不用于 benchmark）

text

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 multi_processor_count=32 wave_size=32
RESULT operator=vector-add implementation=hip-v0 runtime=hip shape=16777216 dtype=float32 block=256 grid=65536 warmup=5 repeat=10 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.368002 median_ms=0.368702 mean_ms=1.549060 effective_bandwidth_gbs=546.041483 max_abs_error=0

Profiler 文件生成记录

text

W0920 08:28:05.453606 1407749 simple_timer.cpp:55] [rocprofv3] tool initialization ::     0.198063 sec
W0920 08:28:05.454094 1407749 simple_timer.cpp:55] [rocprofv3] 'chapter8/results/walkthrough-20260920/build/vector_add_hip --version v0 --size 16777216 --block 256 --warmup 5 --repeat 10 --seed 20260920' ::     0.000000 sec
W0920 08:28:05.459627 1407749 tool.cpp:3348] HSA version 1.21.0 initialized (instance=0)
W0920 08:28:05.553311 1407749 queue_controller.cpp:170] [queue-intercept] device-memory ring-buffer requested but profiling requires system-memory InterceptQueue; falling back to system-memory ring for agent 109707699042784 (priority and CU-mask from the descriptor are not preserved)
W0920 08:28:05.679374 1407749 simple_timer.cpp:55] [rocprofv3] 'chapter8/results/walkthrough-20260920/build/vector_add_hip --version v0 --size 16777216 --block 256 --warmup 5 --repeat 10 --seed 20260920' ::     0.225279 sec
E0920 08:28:05.700510 1407749 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/pmc-waves/baseline_counter_collection.csv
E0920 08:28:05.705534 1407749 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/pmc-waves/baseline_agent_info.csv
W0920 08:28:05.706677 1407749 simple_timer.cpp:55] [rocprofv3] output generation ::     0.015960 sec
W0920 08:28:05.706711 1407749 simple_timer.cpp:55] [rocprofv3] tool finalization ::     0.015996 sec

baseline_counter_collection.csv

text

"Correlation_Id","Dispatch_Id","Agent_Id","Queue_Id","Process_Id","Thread_Id","Grid_Size","Kernel_Id","Kernel_Name","Workgroup_Size","LDS_Block_Size","Scratch_Size","VGPR_Count","Accum_VGPR_Count","SGPR_Count","Counter_Name","Counter_Value","Start_Timestamp","End_Timestamp"
8,8,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GRBM_GUI_ACTIVE",587211.000000,908801672169834,908801672534236
8,8,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"SQ_WAVES_sum",524288.000000,908801672169834,908801672534236
9,9,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GRBM_GUI_ACTIVE",528736.000000,908801684159393,908801684486634
9,9,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"SQ_WAVES_sum",524288.000000,908801684159393,908801684486634
10,10,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GRBM_GUI_ACTIVE",527469.000000,908801684534475,908801684861637
10,10,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"SQ_WAVES_sum",524288.000000,908801684534475,908801684861637
11,11,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GRBM_GUI_ACTIVE",528474.000000,908801684908317,908801685235639
11,11,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"SQ_WAVES_sum",524288.000000,908801684908317,908801685235639
12,12,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GRBM_GUI_ACTIVE",529355.000000,908801685282319,908801685609441
12,12,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"SQ_WAVES_sum",524288.000000,908801685282319,908801685609441
13,13,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GRBM_GUI_ACTIVE",529368.000000,908801685655721,908801685982803
13,13,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"SQ_WAVES_sum",524288.000000,908801685655721,908801685982803
14,14,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GRBM_GUI_ACTIVE",529518.000000,908801686028883,908801686356125
14,14,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"SQ_WAVES_sum",524288.000000,908801686028883,908801686356125
15,15,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GRBM_GUI_ACTIVE",528427.000000,908801686402365,908801686729287
15,15,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"SQ_WAVES_sum",524288.000000,908801686402365,908801686729287
16,16,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GRBM_GUI_ACTIVE",530411.000000,908801686775168,908801687103049
16,16,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"SQ_WAVES_sum",524288.000000,908801686775168,908801687103049
17,17,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GRBM_GUI_ACTIVE",529243.000000,908801687149930,908801687477012
17,17,"Agent 1",2,1407749,1407749,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"SQ_WAVES_sum",524288.000000,908801687149930,908801687477012

读取 baseline_counter_collection.csv 时，先找 Counter_Name=SQ_WAVES_sum。同一个 Dispatch_Id 下有不同计数器的记录，不要把 CSV 行数当作 kernel 次数。本次每个目标 dispatch 的值都是 524,288。本例实际采用 wave32，可以按分工核对：

65, 536 blocks \times \frac{256 threads}{32 threads/wave} = 524, 288 waves

这说明本次 wave 计数与已知启动规模吻合。它并没有告诉我们显存传输了多少字节，也不表示 wave 数越少就越快。

我们还试了 GL2C_HIT_sum、GL2C_MISS_sum 与 FetchSize，本机当前电源状态下得到的记录均为零，没有取得足以解释缓存命中率或物理流量的有效证据。AMD 的 SDK 限制说明要求 gfx11/gfx12 计数器采集使用稳定 power state；本次保留机器原来的 AUTO 状态。因此本章不据这些零值计算命中率，也不把地址反例的慢直接换算成 DRAM 事务数。

查看本次缓存与流量计数器的尝试：命令及原始零值记录

两组分别采集，程序与参数不变；结果保存在不同目录。

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
mkdir "$RUN/pmc-l2" && \
rocprofv3 --pmc GL2C_HIT_sum GL2C_MISS_sum \
  --kernel-include-regex vector_add_v0 --kernel-iteration-range 7-16 \
  --output-format csv --output-directory "$RUN/pmc-l2" \
  --output-file baseline -- \
  "$RUN/build/vector_add_hip" --version v0 --size 16777216 --block 256 \
  --warmup 5 --repeat 10 --seed 20260920

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
mkdir "$RUN/pmc-fetch" && \
rocprofv3 --pmc FetchSize \
  --kernel-include-regex vector_add_v0 --kernel-iteration-range 7-16 \
  --output-format csv --output-directory "$RUN/pmc-fetch" \
  --output-file baseline -- \
  "$RUN/build/vector_add_hip" --version v0 --size 16777216 --block 256 \
  --warmup 5 --repeat 10 --seed 20260920

缓存计数器：采集进程输出

text

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 multi_processor_count=32 wave_size=32
RESULT operator=vector-add implementation=hip-v0 runtime=hip shape=16777216 dtype=float32 block=256 grid=65536 warmup=5 repeat=10 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.381122 median_ms=0.381562 mean_ms=0.421637 effective_bandwidth_gbs=527.638639 max_abs_error=0

缓存计数器：原始 CSV

text

"Correlation_Id","Dispatch_Id","Agent_Id","Queue_Id","Process_Id","Thread_Id","Grid_Size","Kernel_Id","Kernel_Name","Workgroup_Size","LDS_Block_Size","Scratch_Size","VGPR_Count","Accum_VGPR_Count","SGPR_Count","Counter_Name","Counter_Value","Start_Timestamp","End_Timestamp"
8,8,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_HIT_sum",0.00000000e+00,908802309333301,908802309697943
8,8,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_MISS_sum",0.00000000e+00,908802309333301,908802309697943
9,9,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_HIT_sum",0.00000000e+00,908802309758303,908802310263826
9,9,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_MISS_sum",0.00000000e+00,908802309758303,908802310263826
10,10,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_HIT_sum",0.00000000e+00,908802310323706,908802310650747
10,10,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_MISS_sum",0.00000000e+00,908802310323706,908802310650747
11,11,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_HIT_sum",0.00000000e+00,908802310710388,908802311037749
11,11,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_MISS_sum",0.00000000e+00,908802310710388,908802311037749
12,12,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_HIT_sum",0.00000000e+00,908802311097390,908802311424631
12,12,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_MISS_sum",0.00000000e+00,908802311097390,908802311424631
13,13,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_HIT_sum",0.00000000e+00,908802311484351,908802311811633
13,13,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_MISS_sum",0.00000000e+00,908802311484351,908802311811633
14,14,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_HIT_sum",0.00000000e+00,908802311870913,908802312197915
14,14,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_MISS_sum",0.00000000e+00,908802311870913,908802312197915
15,15,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_HIT_sum",0.00000000e+00,908802312256915,908802312584156
15,15,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_MISS_sum",0.00000000e+00,908802312256915,908802312584156
16,16,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_HIT_sum",0.00000000e+00,908802312643517,908802312970678
16,16,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_MISS_sum",0.00000000e+00,908802312643517,908802312970678
17,17,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_HIT_sum",0.00000000e+00,908802313030039,908802313356960
17,17,"Agent 1",2,1407761,1407761,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"GL2C_MISS_sum",0.00000000e+00,908802313030039,908802313356960

FetchSize：采集进程输出

text

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 multi_processor_count=32 wave_size=32
RESULT operator=vector-add implementation=hip-v0 runtime=hip shape=16777216 dtype=float32 block=256 grid=65536 warmup=5 repeat=10 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.399335 median_ms=0.400135 mean_ms=1.500717 effective_bandwidth_gbs=503.147292 max_abs_error=0

FetchSize：原始 CSV

text

"Correlation_Id","Dispatch_Id","Agent_Id","Queue_Id","Process_Id","Thread_Id","Grid_Size","Kernel_Id","Kernel_Name","Workgroup_Size","LDS_Block_Size","Scratch_Size","VGPR_Count","Accum_VGPR_Count","SGPR_Count","Counter_Name","Counter_Value","Start_Timestamp","End_Timestamp"
8,8,"Agent 1",2,1407771,1407771,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"FetchSize",0.00000000e+00,908802923101432,908802923429234
9,9,"Agent 1",2,1407771,1407771,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"FetchSize",0.00000000e+00,908802934145686,908802934473258
10,10,"Agent 1",2,1407771,1407771,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"FetchSize",0.00000000e+00,908802934563220,908802934890512
11,11,"Agent 1",2,1407771,1407771,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"FetchSize",0.00000000e+00,908802934968195,908802935295607
12,12,"Agent 1",2,1407771,1407771,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"FetchSize",0.00000000e+00,908802935373409,908802935700741
13,13,"Agent 1",2,1407771,1407771,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"FetchSize",0.00000000e+00,908802935778384,908802936106076
14,14,"Agent 1",2,1407771,1407771,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"FetchSize",0.00000000e+00,908802936183559,908802936510691
15,15,"Agent 1",2,1407771,1407771,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"FetchSize",0.00000000e+00,908802936588494,908802936915826
16,16,"Agent 1",2,1407771,1407771,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"FetchSize",0.00000000e+00,908802936993149,908802937320601
17,17,"Agent 1",2,1407771,1407771,16777216,23,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)",256,0,0,8,0,128,"FetchSize",0.00000000e+00,908802937397804,908802937725055

这里核对的是目标 dispatch 的 Counter_Value，不是上方采集进程打印的 event 耗时。全零尚不能解释原来的快慢，下一步应先核对采集条件，而不是用这些零值计算命中率或带宽。

一行加法对应哪些指令？ 下例跟踪 baseline v0 的一个局部硬件位置，将源码对应到实际加载、等待、计算与写回。它用于理解指令报告，不用来解释 v3 或把局部周期换算成整卡瓶颈。

选读：用 ATT 查看 baseline 的源码与指令对应

若要观察某行源码对应哪些指令，可以用局部 wave 跟踪补充证据。下面只演示 baseline 的源码对应，不归因某个优化版本的时间差。SDK 的 ATT 支持表已列出 gfx1201，模式为 trace-only；本例不使用 Instinct 示例里的 --att-perfcounters。

这是进阶选读，不影响前面完成一轮优化实验。本章提供一个短脚本，仍使用完整的 vector_add_hip.hip：单独编译带行号的诊断程序，选择 v0 的第 7 次调用，即预检与 5 次预热之后的一次调用，再限定一个 shader engine 与 SIMD。

完整采集脚本：code/part2-kernels/chapter8/profile_att.sh

bash

#!/usr/bin/env bash
# Run from code/part2-kernels after source ./activate-rocm.sh.
set -euo pipefail
if [[ $# != 1 || -z "$1" ]]; then
  echo 'Usage: bash chapter8/profile_att.sh NEW_OUTPUT_DIRECTORY' >&2
  exit 2
fi
att_dir=$1
if [[ -e "$att_dir" ]]; then
  echo "Output already exists: $att_dir" >&2
  exit 2
fi
[[ -f chapter8/vector_add_hip.hip ]] || { echo 'Run from code/part2-kernels.' >&2; exit 2; }
mapfile -t sdk_libs < <(python - <<'PY'
from pathlib import Path
import sysconfig
site = Path(sysconfig.get_path('purelib'))
for package in ('_rocm_sdk_devel', '_rocm_sdk_core'):
    print(site / package / 'lib')
PY
)
[[ ${#sdk_libs[@]} == 2 && -d "${sdk_libs[0]}" && -d "${sdk_libs[1]}" ]] || {
  echo 'Activate the ROCm 10.0 environment first.' >&2; exit 2;
}
mkdir -p "$att_dir/build" "$att_dir/trace"
export LD_LIBRARY_PATH="${sdk_libs[0]}:${sdk_libs[1]}${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}"
exec 3>"$att_dir/commands.log"
export BASH_XTRACEFD=3
set -x
hipcc --offload-arch=gfx1201 -O3 -std=c++17 -gline-tables-only \
  chapter8/vector_add_hip.hip -o "$att_dir/build/vector_add_hip_debug" \
  >"$att_dir/compile.stdout.log" 2>"$att_dir/compile.stderr.log"
rocprofv3 --att --att-library-path "${sdk_libs[@]}" \
  --att-simd-select 0 --att-shader-engine-mask 0x1 \
  --att-buffer-size 67108864 --kernel-include-regex vector_add_v0 \
  --kernel-iteration-range 7 --output-format csv \
  --output-directory "$att_dir/trace" --output-file baseline -- \
  "$att_dir/build/vector_add_hip_debug" --version v0 --size 16777216 \
  --block 256 --warmup 5 --repeat 1 --seed 20260920 \
  >"$att_dir/profile.stdout.log" 2>"$att_dir/profile.stderr.log"
python chapter8/validate_att.py "$att_dir/trace" \
  >"$att_dir/validation.stdout.log" 2>"$att_dir/validation.stderr.log" || {
    cat "$att_dir/validation.stderr.log" >&2; exit 1;
  }
set +x
cat "$att_dir/validation.stdout.log"
echo "ATT captured and decoded: $att_dir/trace"
echo 'One selected dispatch / shader engine / SIMD; not whole-GPU counters.'

脚本最后会调用随附的 validate_att.py，核对目标 dispatch 的指令统计、kernel 名称和解码 JSON；产物缺失或为空时会报错，不会只凭 profiler 正常退出就宣布解码成功。

完整产物检查脚本：code/part2-kernels/chapter8/validate_att.py

python

"""Validate the single-dispatch ATT output expected by profile_att.sh."""
from __future__ import annotations

import argparse
import csv
import json
from pathlib import Path
import re


def validate_att(directory: Path) -> dict:
    candidates = sorted(directory.glob('stats_*.csv'))
    if len(candidates) != 1:
        raise ValueError('expected exactly one decoded stats CSV')
    stats = candidates[0]
    match = re.fullmatch(r'stats_(ui_output_agent_\d+_dispatch_8)\.csv', stats.name)
    if match is None:
        raise ValueError('expected decoded stats for dispatch 8')
    integer_fields = ('CodeObj', 'Vaddr', 'Hitcount', 'Latency', 'Stall', 'Idle')
    required = {*integer_fields, 'Instruction', 'Source'}
    with stats.open(newline='') as file:
        reader = csv.DictReader(file, strict=True)
        fields = reader.fieldnames or []
        if len(fields) != len(set(fields)) or not required <= set(fields):
            raise ValueError('stats CSV has missing or duplicate required columns')
        rows = list(reader)
    if not rows:
        raise ValueError('stats CSV contains no instruction rows')
    for row in rows:
        if None in row or any(value is None for value in row.values()):
            raise ValueError('stats CSV row width does not match its header')
        if not row['Instruction'].strip():
            raise ValueError('stats CSV contains an empty instruction')
        if any(re.fullmatch(r'[0-9]+', row[field]) is None for field in integer_fields):
            raise ValueError('stats CSV counts and addresses must be non-negative integers')
    if not any(re.search(r'\bvector_add_v0\(', row['Source']) for row in rows):
        raise ValueError('stats CSV does not identify vector_add_v0')
    if not any(int(row['Hitcount']) > 0 for row in rows):
        raise ValueError('stats CSV contains no positive instruction Hitcount')
    ui = directory / match[1]
    json_files = sorted(ui.glob('*.json'))
    if not json_files or not (ui / 'code.json').is_file():
        raise ValueError('corresponding UI directory or code.json is missing')
    for path in json_files:
        document = json.loads(path.read_text())
        if not document:
            raise ValueError(f'empty decoded JSON: {path.name}')
    code = json.loads((ui / 'code.json').read_text())
    if not isinstance(code, dict) or not code.get('code'):
        raise ValueError('code.json contains no decoded instructions')
    if not list(ui.glob('*_wv*.json')):
        raise ValueError('UI directory contains no decoded wave JSON')
    return dict(dispatch=8, kernel='vector_add_v0', instruction_rows=len(rows),
                json_files=len(json_files), stats=stats.name, ui=ui.name)


def main() -> int:
    parser = argparse.ArgumentParser(description=__doc__)
    parser.add_argument('directory', type=Path)
    args = parser.parse_args()
    try:
        result = validate_att(args.directory)
    except (OSError, ValueError, csv.Error) as error:
        parser.exit(1, f'ATT validation failed: {error}\n')
    print('ATT_VALIDATED ' + ' '.join(f'{key}={value}' for key, value in result.items()))
    return 0


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

bash

test -d "${RUN:?请先完成环境与结果目录准备}/build" && \
bash chapter8/profile_att.sh "$RUN/att"

查看 ATT 运行结果、采集记录与解码后的指令统计

脚本的终端输出

text

ATT_VALIDATED dispatch=8 kernel=vector_add_v0 instruction_rows=29 json_files=4134 stats=stats_ui_output_agent_3208_dispatch_8.csv ui=ui_output_agent_3208_dispatch_8
ATT captured and decoded: chapter8/results/walkthrough-20260920/att-verified/trace
One selected dispatch / shader engine / SIMD; not whole-GPU counters.

采集进程输出（不用于 benchmark）

text

ENV runtime=hip gpu="AMD Radeon RX 9070 XT" arch=gfx1201 multi_processor_count=32 wave_size=32
RESULT operator=vector-add implementation=hip-v0 runtime=hip shape=16777216 dtype=float32 block=256 grid=65536 warmup=5 repeat=1 seed=20260920 timed=1 correct=OK precheck=OK postcheck=OK min_ms=0.389084 median_ms=0.389084 mean_ms=0.389084 effective_bandwidth_gbs=517.437330 max_abs_error=0

Profiler 记录

text

W0920 08:41:31.355825 1414466 simple_timer.cpp:55] [rocprofv3] tool initialization ::     0.065424 sec
W0920 08:41:31.356418 1414466 simple_timer.cpp:55] [rocprofv3] 'chapter8/results/walkthrough-20260920/att-verified/build/vector_add_hip_debug --version v0 --size 16777216 --block 256 --warmup 5 --repeat 1 --seed 20260920' ::     0.000000 sec
W0920 08:41:31.365662 1414466 tool.cpp:3348] HSA version 1.21.0 initialized (instance=0)
W0920 08:41:31.460928 1414466 queue_controller.cpp:170] [queue-intercept] device-memory ring-buffer requested but profiling requires system-memory InterceptQueue; falling back to system-memory ring for agent 107007324322592 (priority and CU-mask from the descriptor are not preserved)
E0920 08:41:31.478486 1414466 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/att-verified/trace/baseline_gfx1201_code_object_id_1.out
E0920 08:41:31.491417 1414466 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/att-verified/trace/baseline_gfx1201_code_object_id_2.out
E0920 08:41:31.537903 1414468 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/att-verified/trace/baseline_3208_shader_engine_0_8.att
W0920 08:41:31.566081 1414466 simple_timer.cpp:55] [rocprofv3] 'chapter8/results/walkthrough-20260920/att-verified/build/vector_add_hip_debug --version v0 --size 16777216 --block 256 --warmup 5 --repeat 1 --seed 20260920' ::     0.209663 sec
E0920 08:41:31.579793 1414466 output_stream.cpp:110] Opened result file: chapter8/results/walkthrough-20260920/att-verified/trace/baseline_agent_info.csv
W0920 08:41:32.202873 1414466 simple_timer.cpp:55] [rocprofv3] output generation ::     0.623269 sec
W0920 08:41:32.202914 1414466 simple_timer.cpp:55] [rocprofv3] tool finalization ::     0.623313 sec

stats_ui_output_agent_3208_dispatch_8.csv 的归档副本

text

"CodeObj","Vaddr","Instruction","Hitcount","Latency","Stall","Idle","Source"
2,13568,"; _ZN12_GLOBAL__N_113vector_add_v0EPKfS1_Pfm",0,0,0,0,"(anonymous namespace)::vector_add_v0(float const*, float const*, float*, unsigned long)"
2,13568,"s_clause 0x1",4124,4124,0,131694,"<environment-part-root>/.venv/lib/python3.12/site-packages/_rocm_sdk_devel/include/hip/amd_detail/amd_hip_runtime.h:258 -> <environment-part-root>/.venv/lib/python3.12/site-packages/_rocm_sdk_devel/include/hip/amd_detail/amd_hip_runtime.h:317 -> <walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:40"
2,13572,"s_load_b32 s8, s[0:1], 0x2c",4124,11020,6896,0,"<environment-part-root>/.venv/lib/python3.12/site-packages/_rocm_sdk_devel/include/hip/amd_detail/amd_hip_runtime.h:258 -> <environment-part-root>/.venv/lib/python3.12/site-packages/_rocm_sdk_devel/include/hip/amd_detail/amd_hip_runtime.h:317 -> <walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:40"
2,13580,"s_load_b256 s[0:7], s[0:1], 0x0",4124,4143,19,3,"<environment-part-root>/.venv/lib/python3.12/site-packages/_rocm_sdk_devel/include/hip/amd_detail/amd_hip_runtime.h:258 -> <environment-part-root>/.venv/lib/python3.12/site-packages/_rocm_sdk_devel/include/hip/amd_detail/amd_hip_runtime.h:317 -> <walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:40"
2,13588,"v_mov_b32_e32 v1, 0",4124,4124,0,1806,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:40"
2,13592,"s_wait_kmcnt 0x0",4124,159748,155624,0,"<environment-part-root>/.venv/lib/python3.12/site-packages/_rocm_sdk_devel/include/hip/amd_detail/amd_hip_runtime.h:258 -> <environment-part-root>/.venv/lib/python3.12/site-packages/_rocm_sdk_devel/include/hip/amd_detail/amd_hip_runtime.h:317 -> <walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:40"
2,13596,"s_and_b32 s8, s8, 0xffff",4124,4124,0,47,"<environment-part-root>/.venv/lib/python3.12/site-packages/_rocm_sdk_devel/include/hip/amd_detail/amd_hip_runtime.h:258 -> <environment-part-root>/.venv/lib/python3.12/site-packages/_rocm_sdk_devel/include/hip/amd_detail/amd_hip_runtime.h:317 -> <walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:40"
2,13604,"s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)",0,0,0,0,"<environment-part-root>/.venv/lib/python3.12/site-packages/_rocm_sdk_devel/include/hip/amd_detail/amd_hip_runtime.h:258 -> <environment-part-root>/.venv/lib/python3.12/site-packages/_rocm_sdk_devel/include/hip/amd_detail/amd_hip_runtime.h:317 -> <walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:40"
2,13608,"v_mad_co_u64_u32 v[0:1], null, s8, ttmp9, v[0:1]",4124,16496,0,34124,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:40"
2,13616,"s_delay_alu instid0(VALU_DEP_1)",0,0,0,0,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:40"
2,13620,"v_cmp_gt_u64_e32 vcc_lo, s[6:7], v[0:1]",4124,8248,0,17131,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:41"
2,13624,"s_and_saveexec_b32 s6, vcc_lo",4124,4124,0,62166,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:41"
2,13628,"s_cbranch_execz 27",4124,4124,0,32920,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:41"
2,13632,"v_lshlrev_b64_e32 v[0:1], 2, v[0:1]",4124,8248,0,9557,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:?"
2,13636,"s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)",0,0,0,0,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:?"
2,13640,"v_add_co_u32 v2, vcc_lo, s0, v0",4124,4124,0,16613,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:?"
2,13648,"v_add_co_ci_u32_e64 v3, null, s1, v1, vcc_lo",4124,5267,1143,16446,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:?"
2,13656,"v_add_co_u32 v4, vcc_lo, s2, v0",4124,4124,0,71,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:?"
2,13664,"s_wait_alu depctr_va_vcc(0)",0,0,0,0,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:?"
2,13668,"v_add_co_ci_u32_e64 v5, null, s3, v1, vcc_lo",4124,4124,0,66921,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:?"
2,13676,"global_load_b32 v2, v[2:3], off",4124,8255,7,61975,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:42"
2,13688,"global_load_b32 v3, v[4:5], off",4124,8248,0,0,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:42"
2,13700,"v_add_co_u32 v0, vcc_lo, s4, v0",4124,4124,0,34588,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:42"
2,13708,"s_wait_alu depctr_va_vcc(0)",0,0,0,0,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:42"
2,13712,"v_add_co_ci_u32_e64 v1, null, s5, v1, vcc_lo",4124,4124,0,62432,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:42"
2,13720,"s_wait_loadcnt 0x0",4124,11111887,11107763,0,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:42"
2,13724,"v_add_f32_e32 v2, v2, v3",4124,4124,0,124,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:42"
2,13728,"global_store_b32 v[0:1], v2, off",4124,12372,0,65950,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:42"
2,13740,"s_endpgm",4124,1481481,0,0,"<walkthrough-root>/code/part2-kernels/chapter8/vector_add_hip.hip:44"

本次生成原始 .att、对应 code object，以及解码后的 stats_ui_output_agent_3208_dispatch_8.csv 与 ui_output_agent_3208_dispatch_8/；文件名中的 agent 编号可能随运行变化。终端中的 ATT_VALIDATED 确认目标为 dispatch 8、kernel 为 vector_add_v0，并检查了 29 行指令统计和 4,134 个非空的有效 JSON 文件。输出记录中只有选定的一次 dispatch。脚本显式指定当前 ROCm wheel 的解码库目录，是因为本次直接采用默认搜索路径时出现了 Error loading decoder: 37；指定路径后采集与解码均成功。脚本会保留实际命令、编译日志、采集日志与产物检查结果。

先读 Instruction 与 Source 两列。本次 vector_add_hip.hip:42 的一行加法，映射出两条 global_load_b32，随后是 s_wait_loadcnt、v_add_f32_e32 与 global_store_b32。它把源码中的“读两个数、等待数据、相加、写回”对应到了真实指令。

再看 Hitcount：上述指令在本次局部跟踪中各记录了 4,124 次。它是对已追踪 wave 累加的该指令执行次数；本例每个被跟踪的 wave 执行这些指令各一次。它和上一节全 dispatch 的 524,288 个 wave 不是同一范围。Latency、Stall 与 Idle 是周期统计，不是 ms；即使等待指令对应较大数值，也需要结合跟踪范围、依赖和受控对照分析，不能直接宣布它贡献了整个 kernel 的多少百分比。字段定义见 ATT Stats CSV。

这些轨迹只代表选定硬件位置与这一次调用。局部 wave 的指令时间不能直接加成整个 GPU 的 kernel 时间，也不能拿启用 ATT 后的 event 耗时替换 benchmark。解码目录可交给 ROCprof Compute Viewer继续查看；本章验证到采集与解码，尚未验证该查看器在本机上的图形展示。

8.6 HIP 与 Triton 对照

8.6.1 两条路线调整的对象

问题	HIP 直接表达	Triton 直接表达
一份工作负责哪些数据	线程下标、循环步长、向量分组	program 的 offsets 与逻辑 tile
尾部怎样保护	标量 `if` 与标量尾部循环	load/store 各自的 mask
本章改变的执行粒度	固定 kernel 后调整 grid	固定 `num_warps` 后调整 `BLOCK_SIZE`
谁完成底层映射	源码指定线程到地址的关系，编译器生成指令	编译器将 tile 分配给 lane、寄存器与执行组

HIP 的 thread 与 Triton 的 program 不是一一对应关系。比较两种实现时，先核对负责哪些输出、读哪些输入、怎样处理边界，再核对时间与资源，而不是把两种语言的参数名称直接互换。

8.6.2 将每轮结果放回原始基线

HIP 与 Triton 向量加法各版本的完整时间汇总

图 8.9　相同主输入下的代表配置总览；下表列出全部 11 个配置。地址对照是独立诊断，不是从 v0 开始的线性升级链；这张图用于回看各轮选择，不替代前面的受控比较。

配置	Median（ms）	三进程范围（ms）
`hip-v0`	0.334861	0.333301–0.349441
`hip-v1-contiguous`	0.369681	0.366240–0.370081
`hip-v1-strided`	2.538206	2.505867–2.576887
`hip-v2-g65536`	0.350301	0.332861–0.355121
`hip-v2-g2048`	0.357582	0.353582–0.361721
`hip-v2-g256`	0.343640	0.334301–0.344421
`hip-v3-g256`	0.342121	0.342002–0.342500
`triton-t0-b256`	0.334061	0.333781–0.349521
`triton-t1-b512`	0.335441	0.335401–0.340481
`triton-t1-b1024`	0.337701	0.336301–0.338220
`triton-t1-b2048`	0.341241	0.340961–0.343221

hip-v0 与 triton-t0-b256 的三进程测量范围高度重叠，基于当前严谨的统计数据，我们不能得出“某种语言在性能上压倒另一种”的轻率结论。更为客观的认识是：在当前的大规模输入下，两条路线最自然的基线实现均已展现出极高的访存效率；后续试图通过更复杂的技巧进行重构时，唯有拿出同口径下更显著的提速证据，才值得替换掉简洁明了的基线。

正文引用的单进程输出、三进程统计与五张对比图表，均完整归档于配套代码库的 evidence/walkthrough/ 目录中，确保所有实验数据源头均可独立追溯与复核。

8.7 负结果、适用边界与下一步

一次改动没有提速，也能帮助我们判断在当前条件下是否值得保留。记录观察结果时，要同时说明它还不能解释什么：

观察到的实测现象	据此可以作出的工程决定	依然不能直接推论的假设
受控地址排列对比出现 6.87 倍断崖落差	坚决保留同轮连续访存模式，在后续所有算子中持续审视访问连续性	不能脱离硬件计数器直接将理论分桶等同于真实 DRAM 事务数
Grid 或 Tile 缩放未见稳定收益	保留最初简洁的配置，不盲目为了减少线程/Program 而加重单任务负担	不能推广认为“更少的 Block/Program 在所有场景下都普遍低效”
引入 `float4` 向量化后无稳定提速	保留原生标量加法写法，避免徒增代码复杂度	不能断言“向量类型在硬件层面上绝对没有生成更宽的访存指令”

本章得出的最佳配置结论严格受限于当前的大规模输入。当计算场景切换到小尺寸张量或非连续内存布局时，应当重新运行基线并评估候选版本，切忌将大数组下的经验生搬硬套。

如果未来我们需要实现 ReLU、GELU 或 Scale & Bias，由于输出元素依然保持严格独立，本章建立的网格划分与边界保护策略均可直接迁移复用。如果考虑将 Add 与后续的激活算子进行算子融合，则必须严格测量“两次独立算子调用”与“单次融合内核调用”的完整端到端路径，相关的融合收益与代价我们将在第 12 章深入探讨。

8.8 复跑与练习

8.8.1 复现逐轮实验

在仓库根目录进入本篇环境：

bash

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

单项命令帮助你理解一个改动。要比较运行波动，用批量入口自动运行 11 个配置，每个配置启动 3 个独立进程，并轮换运行顺序。这里显式选择 main：做小边界检查与 benchmark，不采集 profiler。输出目录必须尚不存在：

bash

python chapter8/run_rounds.py --output chapter8/results/my-run --phase main

程序默认同样为 main。它在测量前检查 1、31、32、33、255、256、257、1027 等小长度；主输入每个进程仍执行前后检查。profile 与 validation 是额外阶段，不是每次实验必做的下一步；profile 会采集全部配置，当前章节通常只需 8.4 中针对问题的单项命令。已有主实验目录不会被覆盖。

若需要在独立进程中再次确认选定配置，并检查各 tile 的专属边界，可以执行下面的补充验收。这是单独的验收步骤，不包含在上面的默认运行中：

bash

python chapter8/verify_rounds.py \
  --frozen-run chapter8/results/my-run \
  --output chapter8/results/my-confirmation

确认步骤复用 my-run 中保存的 kernel，结果写进新目录；不覆盖原扫描。8.4 的单项命令用于看清每一步，这里的运行器用于自动保存完整的多进程矩阵。

从自己的汇总重画本章五类比较图：

bash

python chapter8/plot_rounds.py \
  --summary chapter8/results/my-run/summary/summary.csv \
  --manifest chapter8/results/my-run/summary/manifest.json \
  --out-dir chapter8/results/my-run/figures

运行目录中，samples/ 保存每次 event 时间，commands.jsonl 保存实际命令，logs/ 保存环境、程序输出和检查结果，profiles/ 保存原始 trace，source/ 保存该次源码，summary/ 保存汇总与环境清单。仅执行 main 时，profiles/ 为空是正常现象。上面的绘图命令显式读取你自己的汇总；不要省略数据参数，否则脚本的默认值会读取另一组归档。正文五张图对应 evidence/walkthrough/；evidence/rounds/ 是先前完整运行器实验的独立快照，两组不混算成绩。

采集日期、源码身份与详细条件保存在每组实验记录中。新增结果以自己的运行目录保存，原始样本、环境、源码哈希和实际参数一起保留；不同日期、shape 或计时区间的结果不混成一个中位数。

8.8.2 练习与验收

选择一个新的非整除长度，先预测最后一个 HIP block 和 Triton program 中哪些位置有效，再检查输出。正确性通过后才计时。
固定 v2 的源码与 block size，增加一个 grid 候选。把结果与现有 grid 配置放在同一张图中，说明这次是否值得保留。
固定 Triton 的 num_warps=4，改变一个 tile 候选；记录 program 数和实测范围。若差异很小，说明为什么还不能宣布胜者。
将算术替换为 output[i] = alpha * input[i] + beta。重新计算逻辑字节与 FLOP，先写出改动假设，再校验和比较；如果需要 profiling，写明它要回答的问题与预期字段。

完成时应能独立回答：我改了哪项分工，为什么值得试，正确性怎样确认，比较图支持什么决定，还有什么原因没有证实。如果决定进一步 profiling，还要说明现有测量缺少什么证据、工具能否取得它，以及不同结果会怎样改变下一步。某个优化一定成功，不是验收条件。

本章小结

输出独立，分工可以改变：一个输出只依赖对应输入，可以比较不同的线程分工与数据块大小。低算术强度提示优先检查数据访问，但不能据此断定显存带宽是唯一瓶颈。
两条路线各有控制方式：HIP 从线程索引出发，比较地址排列、grid-stride 循环与四元素分组；Triton 从分块张量出发，比较 tile 大小。
每轮都检查并复测：先核对结果，再按相同口径测量，与基线及上一轮比较。独立复测与样本范围帮助判断收益是否稳定，没有提速的结果也要保留。
有具体问题时再做 profiling：正确性与性能比较可以支持保留或回退的决定；需要进一步解释时，再选择能够提供相关证据的工具。

接下来的第 9 章：Reduction 归约算子研究一个输出依赖多个输入的情况：怎样组织线程协作，并把各组的局部结果合并为最终答案。

第8章 08Element-Wise：逐元素算子 ​

本章导读 ​

8.1 逐元素算子的数据依赖 ​

8.1.1 输出之间的独立性 ​

8.1.2 HIP 与 Triton 的实现视角 ​

8.1.3 与归约算子的区别 ​

8.2 固定数学语义与正确性标准 ​

8.2.1 四元素示例 ​

8.2.2 参考实现与检查条件 ​

8.2.3 非整除长度的边界 ​

8.3 建立成本模型和瓶颈假设 ​

8.3.1 逻辑计算量与数据量 ​

8.3.2 有效带宽怎样算 ​

8.4 实现向量加法 ​

8.4.1 HIP 基线：一个线程处理一个元素 ​

8.4.2 地址对照：只改变同轮访问顺序 ​

8.4.3 Grid-Stride：减少 block 数是否有收益 ​

8.4.4 float4：一次表达四个元素是否更快 ​

8.4.5 HIP 路线的实验决策 ​

8.4.6 Triton 基线：一个 program 处理一段元素 ​

8.4.7 Tile 扫描：固定执行参数，只改数据粒度 ​

8.4.8 从手工比较到自动调优 ​

8.5 按问题选择性能分析 ​

8.5.1 从问题到字段 ​

8.5.2 时间线与工具适用范围 ​

8.5.3 进阶诊断：计数器与指令跟踪 ​

8.6 HIP 与 Triton 对照 ​

8.6.1 两条路线调整的对象 ​

8.6.2 将每轮结果放回原始基线 ​

8.7 负结果、适用边界与下一步 ​

8.8 复跑与练习 ​

8.8.1 复现逐轮实验 ​

8.8.2 练习与验收 ​

本章小结 ​

延伸阅读 ​

第8章 Element-Wise：逐元素算子

本章导读

8.1 逐元素算子的数据依赖

8.1.1 输出之间的独立性

8.1.2 HIP 与 Triton 的实现视角

8.1.3 与归约算子的区别

8.2 固定数学语义与正确性标准

8.2.1 四元素示例

8.2.2 参考实现与检查条件

8.2.3 非整除长度的边界

8.3 建立成本模型和瓶颈假设

8.3.1 逻辑计算量与数据量

8.3.2 有效带宽怎样算

8.4 实现向量加法

8.4.1 HIP 基线：一个线程处理一个元素

8.4.2 地址对照：只改变同轮访问顺序

8.4.3 Grid-Stride：减少 block 数是否有收益

8.4.4 `float4`：一次表达四个元素是否更快

8.4.5 HIP 路线的实验决策

8.4.6 Triton 基线：一个 program 处理一段元素

8.4.7 Tile 扫描：固定执行参数，只改数据粒度

8.4.8 从手工比较到自动调优

8.5 按问题选择性能分析

8.5.1 从问题到字段

8.5.2 时间线与工具适用范围

8.5.3 进阶诊断：计数器与指令跟踪

8.6 HIP 与 Triton 对照

8.6.1 两条路线调整的对象

8.6.2 将每轮结果放回原始基线

8.7 负结果、适用边界与下一步

8.8 复跑与练习

8.8.1 复现逐轮实验

8.8.2 练习与验收

本章小结

延伸阅读