算力手记C++ & COMPUTING SYSTEMS
← AMD 专题与八个进阶单元

AMD / LAB 02 OF 08

一次完整 HIP 调用：让 8 从设备回到 host

从一个线程的正确 kernel 开始，跟踪分配、拷贝、提交、完成和释放。

本单元 4 小时。本公司八个单元共 32h，三家公司共额外 96h 核心实践；这笔时间在共同路线之外，未压入厂商入门的44h。

本页是实践讲义。打开、阅读或复制命令不代表完成。新单元的 SDK、simulator 与硬件命令均未在教材建设中执行；自己的结果应按实际后端保存。

先备与马上可用的例子

13 · 真实资源与RAII：独占拥有33 · Host/device、传输与异步生命周期37 · HIP：从 CPU 数据到一次 kernel 调用

先完成：amd-01 →

直接打开机制图，边看边手推

PREDICT FIRST

具体问题与独立预期

输入与约定

X=[3,-1,4,0,2]，一个 block、一个线程串行累加。

先根据输入写下自己的结果与理由，再展开参考推演；随后运行或检查实现。

对照参考结果与原因

正常回读为 8。将 X[0] 改成 9 后预期 14，必须同步修改测试 oracle。

先用最小串行 kernel 把算术与运行链路分开。它不是性能实现，却容易定位失败发生在哪一层。

机制怎样连接起来

用 host 数组装输入，用 device 分配保存设备可读的数据。host 指针值本身不能替代一份 device 分配。按照本页生命周期表，每个资源都有拥有者、最后一次使用和释放位置，避免错误路径泄漏。

下面源码是依据公开 HIP API 独立写的教学起点，尚未用 HIP 编译或运行。它逐步检查 API 状态：launch 检查和等待完成承担不同作用。同步之后的回读值才可与 oracle 比较。为把本单元范围压小，长度固定为 5；通用长度和并行化留给下一单元。

无设备时运行 33-a/33-b，另写一张 allocated→copied→submitted→completed→read-back 的状态表，并让非法 read-back 返回错误。这个状态机是 CPU 教学工具，不执行 HIP，也不能发现真实驱动故障。

CHOOSE ONE EXECUTION PATH

按手头环境推进

CPU 路径可以先完成。额外的 SDK 或设备验证要有自己的运行记录；本单元的四小时不同时要求完成四个分支。安装等待超过本页预算时，留下阻塞条件，继续可做的数学、代码与协议工作。

CPU · 现有机器 · 本单元尚未验收

条件：现有 macOS 或 Linux；C++20 编译器或 Python 3 标准库。无需加速卡。

本分支做什么：运行 33-a/33-b 并实现非法回读检查。

保存什么证据：保存自己的源码、输入、实际 stdout/stderr 和退出码；只证明 CPU 逻辑。

教材初始执行状态：not-executed。这不是对你个人学习进度的判断。

SDK · 编译与环境 · 本单元尚未验收

条件：已有兼容 ROCm/HIP 的 Linux 工具链。

本分支做什么：保存源码并尝试编译；没有设备时只记录编译结果。

保存什么证据：保存实际工具版本与编译命令；仅编译成功不能记成运行通过。

教材初始执行状态：not-executed。这不是对你个人学习进度的判断。

官方 simulator · 本单元尚未验收

条件：只有明确提供并且版本匹配的官方 simulator 才属于此分支。

本分支做什么：没有规定 HIP simulator；状态机仅为 CPU 模型。

保存什么证据：记录 simulator 名称、版本、输入、日志、退出状态；网页或 CPU 模型不算 simulator。

教材初始执行状态：not-executed。这不是对你个人学习进度的判断。

真实设备 · 本单元尚未验收

条件：用户已有并可使用的兼容设备或实验室环境；本单元不要求购买、租用或提交集群作业。

本分支做什么：执行两个输入，等待完成后回读；保存 API 结果和退出码。

保存什么证据：真实设备输出、同步点、设备型号和复现日志单独保存；未执行保持未通过。

教材初始执行状态：not-executed。这不是对你个人学习进度的判断。

READ → BUILD → BREAK → EXPLAIN

四小时，留下一个完整产出

可拆成多个时段。每次停下时保存代码、输入、实际结果和下一步，不用重新读整篇。没有通过当前检查时继续修正，不靠翻页推进。

1. 预测资源状态 · 30 分钟

  1. 手算每次加法：3→2→6→6→8。
  2. 列 dx、dy、host x 和 answer 的分配位置与合法访问者。

这一段的产出：带箭头的生命周期表

2. 读源码与编译 · 40 分钟

  1. 将下面原创源码保存为 hip_sum.cpp；只有匹配 HIP 环境才编译。
  2. CPU 路径先编译 33-a 和 33-b，抄录实际输出，再手推等价状态表。

这一段的产出：编译记录或 CPU 状态表

3. 建立运行闭环 · 90 分钟

  1. 有设备时运行 hip_sum，检查退出码和 8；为每个 API 的错误信息加操作名，保留第一处错误。
  2. 将 x[0] 改为 9，并把末尾成功判定 answer==8 改为 answer==14，再验证。
  3. 无设备时在自己的 CPU 状态机加入 completed 之前禁止 read-back 的检查，实际触发一次失败。

这一段的产出：两个输入的正确结果与一个错误路径

4. 审查退出路径 · 50 分钟

  1. 枚举第二次分配失败、H2D 失败、launch 失败三条路径，检查已经得到的资源如何清理。
  2. 将固定五元素合同写在 README，不让调用者误以为可传任意 N。
  3. 不要通过向硬件发送非法地址制造故障；错误顺序在 CPU 状态机中注入。

这一段的产出：资源退出路径表

5. 解释责任边界 · 30 分钟

  1. 用英文描述从 x 到 answer 的完整过程。
  2. 在验证状态里明确写 CPU-run、HIP-compiled 或 hardware-run 中实际达到的一层。

这一段的产出：一分钟调用链说明

EXPLICIT ENVIRONMENT / EXPLICIT STATUS

命令与可保存的起点

以下代码按各自环境使用，运行前完成对应步骤并核对版本。编译失败后停止，不运行目录中的旧二进制；每次采集日志使用新的运行目录。设备程序仅运行在你已有且可使用的环境中。

1. 原创完整起点：保存为 hip_sum.cpp；未用 HIP 编译

环境：SDK · 编译与环境 · 状态：not-executed。核对官方 API / 工具说明 ↗

cpp · 按说明保存到自己的练习目录
#include <hip/hip_runtime.h>
#include <array>
#include <iostream>
static bool ok(hipError_t e) { if (e==hipSuccess) return true; std::cerr<<hipGetErrorString(e)<<"\n"; return false; }
__global__ void sum_one(const int* x, int* y) { int s=0; for(int i=0;i<5;++i) s+=x[i]; y[0]=s; }
int main() {
  std::array<int,5> x{3,-1,4,0,2}; int answer=0; int *dx=nullptr,*dy=nullptr;
  if(!ok(hipMalloc(reinterpret_cast<void**>(&dx),x.size()*sizeof(int)))) return 1;
  if(!ok(hipMalloc(reinterpret_cast<void**>(&dy),sizeof(answer)))) { hipFree(dx); return 1; }
  bool good=ok(hipMemcpy(dx,x.data(),x.size()*sizeof(int),hipMemcpyHostToDevice));
  if(good) { sum_one<<<1,1>>>(dx,dy); good=ok(hipGetLastError()); }
  if(good) good=ok(hipDeviceSynchronize());
  if(good) good=ok(hipMemcpy(&answer,dy,sizeof(answer),hipMemcpyDeviceToHost));
  bool freed_y=ok(hipFree(dy)),freed_x=ok(hipFree(dx));
  if(!good||!freed_y||!freed_x) return 1;
  std::cout<<answer<<"\n"; return answer==8 ? 0 : 2;
}
2. 在已配置 HIP 的 Linux 环境执行

环境：真实设备 · 状态：not-executed。核对官方 API / 工具说明 ↗

bash · 按说明保存到自己的练习目录
hipcc -std=c++20 hip_sum.cpp -o hip_sum && ./hip_sum

一个必须能解释的错误

错误情境：删除 hipDeviceSynchronize 后，只根据 launch 返回成功就写“kernel 正确”。

为什么会错：提交成功和执行完成不是同一事件；也还没有与 oracle 比较输出。

怎样修复：检查 launch，再在适当同步点检查异步执行，最后回读比较。不能从提交时间推断执行时间。

EVIDENCE BEFORE ADVANCING

验收与面试追问

  • 8 和 14 的输入、检查常量保持一致。
  • 能够指出每条退出路径负责释放哪些分配。
  • 源码被标为未执行起点；自己的实际编译/运行结果另记。

最后留下这几项

输入和预期、自己的源码或明确标为推演的状态表、实际命令/输出/退出码、一个失败与修复、后端与版本、尚未执行的部分。运行已有教学例子与独立完成修改分别记录。

1. 为什么先用单线程？

对照中文要点与英文回答

降低并发变量数量，先验证分配、传输、启动、同步、回读链路；它不用于主张速度。

A single thread isolates the runtime lifecycle from parallel synchronization. It is a correctness baseline, not an optimization.

2. 两次错误检查能互相代替吗？

对照中文要点与英文回答

launch 状态用于定位提交配置等问题；同步处可能暴露异步执行错误，二者覆盖阶段不同。

Launch errors and asynchronous execution errors are observed at different points. I check both and preserve the first failing operation.

3. 改输入后仍然返回非零一定是 kernel 错吗？

对照中文要点与英文回答

也可能忘记更新末尾的预期 8。先独立算新答案 14，再更新测试合同，而不是删除测试。

A stale expected value can fail a correct kernel. I recompute the oracle and update the test, rather than removing it.

这些是依照本单元机制设计的追问，不是公司真题。个人验收仍依据实际代码、测试、测量与口述。

资料与版本核验

  • HIP programming model ↗

    Host programming; hierarchical thread model

    核验日期：2026-09-05。核验时 HIP latest 页标题为 7.15.0；以安装工具链的版本选择器和本地头文件为准。

  • HIP device memory ↗

    Device allocation and memory copies

    核验日期：2026-09-05。HIP latest / 7.15.0 页面；记录实际 runtime、driver 和设备名称。

  • HIP error handling ↗

    API return values and asynchronous error reporting

    核验日期：2026-09-05。HIP latest / 7.15.0 页面；保存 hipcc --version 与错误发生的 API 名称。

  • Using HIPCC ↗

    hipcc --version; hipconfig --full

    核验日期：2026-09-05。HIPCC latest 文档；用 hipcc --version 记录实际编译器，用 hipconfig --full 记录实际工具链。

正文、问题和实践设计依据公开资料独立撰写。官方页面的示例输出是资料中的结果，不是本机运行证据。latest 链接可能变化，请在自己的复现说明中保留实际版本。