HIP：从 CPU 数据到一次 kernel 调用

先把调用拆成可检查的步骤

我们有五个整数 2、4、6、8、10，希望每个数加一。普通 C++ 循环会得到 3、5、7、9、11。写 HIP 时，数学问题没有变，改变的是工作安排：CPU 负责准备与协调，kernel 是让设备执行的函数，多个逻辑线程各处理不同索引。先理解这三件事，再记 API 名字，才能在结果错误时知道应检查哪一段。

Host 是主程序运行的 CPU 一侧，device 是被调用的计算设备。设备指针不是“另一个普通 vector”：分配方式、可访问位置和释放时机由相应 API 约束。某些系统支持统一地址或内存机制，但不能因此假定任意 host 指针都能交给任意 kernel。最初使用明确的 host 输入、device 输入输出、host 回读结果，数据关系更容易检查。

先画所有权，再写调用

最小流程包括选择设备、分配设备空间、复制输入、启动 kernel、等待结果、复制输出、验证和释放。错误处理贯穿每一步。分配失败后不能继续使用未成功取得的地址；某一步失败也不能遗忘此前成功分配的资源。RAII 依然有用，不过资源的生命必须覆盖所有异步使用，析构函数不会自动解决跨设备依赖。

例一用两个独立 vector 表达两个存储区。复制后修改 device_model 不会改变 host_input；只有明确回读才得到更新结果。它只验证数据流，不能模拟 HIP runtime、地址空间、GPU 并行执行或传输开销。把模型输出称为 GPU 实测，会混淆正在检查的命题。

多出来的线程也有合同

设 block 大小为四，五个元素需要两个 block，产生八个逻辑线程。全局索引等于 block 编号乘 block 大小，再加线程编号。索引零到四做有效工作，五到七必须跳过。向下取整只启动一个 block 会漏掉第五个数；不检查边界则会访问数组以外的存储。空输入最好由 host 明确处理，不依赖零大小 launch 的行为。

逐项读 kernel 和 launch（补充，另估25分钟）

先把普通 C++ 函数读法扩展到这一个 HIP 片段。__global__ 标记由 host 启动、在 device 上执行的 kernel；本例返回类型为 void，结果写入 output。blockIdx.x 是当前块的 x 编号，blockDim.x 是每块 x 方向的线程数，threadIdx.x 是当前线程在块内的 x 编号；编号从零开始。这里的点号仍是成员访问，三个名字由 HIP 提供。

dim3 表示三个无符号维度；dim3(4) 是 (4, 1, 1)，省略的 y、z 都是 1。下面注释中的 hipLaunchKernelGGL 按位置读取，不能把相邻的零随意理解成设备号：

位置	本例填写	作用
1	plus_one	要启动的 kernel
2	dim3((n+3)/4)	grid：启动多少个块
3	dim3(4)	block：每块多少个线程
4	0	每块额外动态共享内存的字节数，本例不申请
5	stream	已准备好的执行队列，工作提交到这个 stream
6 起	input, output, n	按 kernel 形参顺序传递的三个实参

**纸上走一次：**令 n=5。grid 为 (2,1,1)，block 为 (4,1,1)；块零覆盖索引 0–3，块一覆盖 4–7。i < n 让索引 4 写出 11，索引 5–7 跳过，最终五个数是 3、5、7、9、11。只把块大小改成 8 并同步修改 grid 的向上取整式，会得到一个块、仍有三个线程跳过。这里的加法式只用于小整数，不能忽略大 n 的溢出检查。

launch 配置没有表达“已经算完”。两个指针须指向该设备可访问且容量足够的存储，输入要先准备完成，输入、输出与相关资源还要存活到设备使用结束。下方是未执行的语法片段；上述数值是索引推演，CPU 示例不能证明 HIP 调用成功。读法核对自 HIP 7.1.52802 的 kernel 调用与 dim3，对应 docs-7.1.1，核对日期 2026-09-08。

// HIP 语法示意：未在 SDK/设备执行，不是完整 host 程序。
__global__ void plus_one(const int* input, int* output, int n) {
    const int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) output[i] = input[i] + 1;
}
// 已分配和准备 input/output；n > 0，且取值满足 int 范围。
// hipLaunchKernelGGL(plus_one, dim3((n+3)/4), dim3(4),
//                    0, stream, input, output, n);

这里使用很小的整数；真实输入还需检查长度转换、字节数乘法和加法是否溢出。四个线程只是索引算例，不是硬件配置建议。block 大小应结合 kernel 资源用量和设备限制选择，不能从示意图推导性能。

调用返回了，究竟证明什么

Host 发起工作后可能先继续执行。启动错误、异步执行错误和数值错误是三个层次：先检查 API 返回值及 launch 状态，再在规定同步点检查执行状态，最后与独立 reference 比较。打印输出指针只说明得到一串地址，不说明设备完成了写入。

本章固定一种输入类型、一个设备、一个显式 stream 和一次变换。后续再加入流水和 profiling。公开 HIP 文档还要匹配安装版本、目标架构和编译参数；本机例子不要求安装 SDK，也不证明任何目标设备支持情况。

// Original CPU teaching model; not a device simulator or benchmark.
#include <iostream>
#include <vector>
#include <stdexcept>

void require(bool ok) { if (!ok) throw std::runtime_error("model check failed"); }

int main() {
 const std::vector<int> host{2,4,6,8,10};
 auto model=host; for(int& x:model) ++x;
 const auto result=model;
 require(host.front()==2 && result==std::vector<int>({3,5,7,9,11}));
 std::cout<<"host="<<host.front()<<"\nresult=";
 for(std::size_t i=0;i<result.size();++i) std::cout<<(i?" ":"")<<result[i];
 std::cout<<'\n';
}

clang++ -std=c++20 -Wall -Wextra -Wpedantic -Werror -pthread 33-a.cpp -o example && ./example

host=2
result=3 5 7 9 11

// Original CPU teaching model; not a device simulator or benchmark.
#include <iostream>
#include <vector>
#include <stdexcept>

void require(bool ok) { if (!ok) throw std::runtime_error("model check failed"); }

int main() {
 constexpr std::size_t n=5, block=4;
 const auto blocks=n/block+(n%block!=0);
 std::vector<int> visits(n,0); std::size_t skipped=0;
 for(std::size_t b=0;b<blocks;++b) for(std::size_t t=0;t<block;++t) {
  const auto i=b*block+t; if(i<n) ++visits[i]; else ++skipped;
 }
 for(int count:visits) require(count==1);
 require(skipped==3);
 std::cout<<"blocks="<<blocks<<" valid="<<n<<" skipped="<<skipped<<'\n';
}

clang++ -std=c++20 -Wall -Wextra -Wpedantic -Werror -pthread 33-b.cpp -o example && ./example

blocks=2 valid=5 skipped=3