接下来基于rocm-systems库源码文件系统剖析 HIP Graph 的概念、数据结构与执行机制。共分 9 章概念与整体架构数据结构节点与图数据结构辅助资源管理生命周期构建与实例化调度执行一经典路径与 Segment 分段调度执行二AQL Packet 捕获与批量派发调度执行三同步计划 SyncPlan高级特性子图、图内存与图更新实战与总结没有跑过达模型应用所以这部分不太硬本章目标讲清楚 HIP Graph 是什么、为什么需要它以及hip_graph_internal.hpp里三大核心类是如何分工协作的。读完本章你会对整个 Graph 子系统建立起一张地图为后续深入各个模块打好基础。1.1 从一个性能问题说起先看一段最普通的 HIP 代码——它是深度学习、科学计算里最常见的模式一轮迭代里连续提交一批小算子。for(intiter0;iter10000;iter){kernelAg,b,0,stream(...);// 提交 1kernelBg,b,0,stream(...);// 提交 2hipMemcpyAsync(...,stream);// 提交 3kernelCg,b,0,stream(...);// 提交 4}每一次kernel或hipMemcpyAsync调用运行时都要做一整套 CPU 端工作参数校验、查找 kernel、构造 launch 配置创建一个amd::Command对象计算依赖、挂上 event wait list组装 AQL packet、写入硬件队列当单个 kernel 很小几微秒时这套 CPU 端组装的开销可能和 kernel 本身的执行时间相当甚至更大。于是 GPU 频繁饿肚子——算得快但 CPU 喂不上。这就是所谓的CPU launch overhead / dispatch bound问题。关键观察是上面这个循环里每一轮的拓扑结构完全一样变的只是数据。既然结构不变为什么每一轮都要重新组装一遍1.2 Graph 的核心思想一次记录多次重放HIP Graph对标 CUDA Graph给出的答案是把做什么和怎么执行彻底分离重复记录阶段把一批操作及依赖构建成 DAG实例化阶段一次性预计算调度/资源/packet执行阶段直接重放可循环上万次记录阶段把一轮迭代里的所有操作kernel、memcpy、memset……及它们之间的依赖关系记录成一张有向无环图DAG。每个操作是一个节点依赖是边。实例化阶段把这张图编译成一个可执行实例在这一步把所有能提前算的东西全部算好——拓扑排序、流分配、kernel 参数池、信号池、甚至 GPU packet 本身。执行阶段重放时几乎没有 CPU 端组装工作直接把预先准备好的 packet 批量丢给硬件。一句话概括用空间预处理和缓存换时间重放性能。这是理解整个 Graph 子系统所有设计的总纲——后面你会看到大量在实例化时预计算、在执行时零成本的手法。1.3 三大核心类蓝图、节点、实例打开hip_graph_internal.hpp最顶上是一组前向声明classGraph;classGraphNode;classGraphExec;typedefGraphNode*Node;整个子系统就是围绕这三个类展开的它们的分工可以用图纸 → 生产线 → 零件来类比类角色类比关键职责GraphNode节点零件一道工序描述单个操作负责CreateCommand/EnqueueCommandsGraph图蓝图图纸组织节点与依赖负责增删节点、拓扑排序、克隆、内存池GraphExec可执行实例生产线按图纸建好的产线预分配 stream/信号/packet负责Run真正执行GraphNode —— 节点基类GraphNode是所有节点的抽象基类。它持有一个节点在图中的全部结构信息classGraphNode:publichipGraphNodeDOTAttribute{protected:hipGraphNodeType type_;// 节点类型std::vectoramd::Command*commands_;// 该节点对应的底层命令std::vectorNodeedges_;// 后继出边std::vectorNodedependencies_;// 前驱入边int32_tstream_id_-1;// 执行时分配到的流int32_tsegment_id_-1;// 分段调度中的段 IDstd::vectoruint8_t*gpuPackets_;// 预捕获的 GPU packet// ...};注意edges_和dependencies_这一对——它们把 DAG 的邻接关系直接存在了节点里。图的连接靠AddEdgeDep这类方法维护voidAddEdgeDep(constNodechildNode){AddEdge(childNode);// this 记录后继childNode-AddDependency(this);// child 记录前驱}每种具体操作对应一个派生类比如GraphKernelNodekernel 启动、GraphMemcpyNode拷贝、GraphMemsetNode填充、ChildGraphNode子图等等。它们通过重写虚函数CreateCommand/EnqueueCommands/SetParams/clone来定制各自行为——这是典型的模板方法 多态设计第 2 章会展开整个节点家族。Graph —— 可修改的蓝图Graph代表一张可以继续编辑的图。它管理所有顶点并提供构建、查询、调度的能力classGraph{std::vectorNodevertices_;// 图中所有节点std::vectorNodetopoOrder_;// 拓扑排序结果hip::MemoryPool*mem_pool_;// 图关联的内存池// ...public:voidAddNode(constNodenode);boolTopologicalOrder(std::vectorNodeTopoOrder);Graph*clone()const;// ...};Graph是活的你可以往里加节点、删节点、连边。它本身不负责执行——执行是GraphExec的事。GraphExec —— 可执行实例GraphExec通过多继承同时是一个Graph它 clone 了原图的内容和一个引用计数对象classGraphExec:publicamd::ReferenceCountedObject,publicGraph{std::unordered_mapint,std::vectorhip::Stream*parallel_streams_;// 预分配的流GraphKernelArgManager*kernArgManager_;// kernel 参数池GraphSignalManager*signalManager_;// HW 信号池std::unordered_mapint,SegmentBatchsegmentBatches_;// 预捕获的 packet 批SyncPlan sync_plan_;// 预计算的同步方案public:hipError_tInit();hipError_tRun(hip::Stream*stream);// ...};对比GraphGraphExec多出来的成员几乎全是预分配 / 预计算的缓存流池、参数池、信号池、packet 、同步计划。这正是 1.2 节空间换时间思想的直接体现——这些昂贵的准备工作只在Init()里做一次之后每次Run()都直接复用。1.4 类之间的关系全景把节点家族和三大核心类放在一起继承关系如下«abstract»GraphNodeCreateCommand()EnqueueCommands()clone()GraphAddNode()TopologicalOrder()clone()GraphExecInit()Run()ChildGraphNodeGraphKernelNodeGraphMemcpyNodeGraphMemsetNodeGraphHostNodeGraphMemAllocNode这里有两个值得注意的设计点GraphExec继承自Graph可执行实例是一种图——它拥有图的全部结构再叠加执行所需的资源。ChildGraphNode同时继承GraphNode和GraphExec一个子图节点既是父图里的一个节点本身又是一张可执行的图。这种双重身份让图可以嵌套是实现层级化图的关键第 8 章详解。1.5 一次完整生命周期鸟瞰最后把用户 API 和内部类串起来看一个 Graph 从生到死的完整流程GraphExecGraph用户代码GraphExecGraph用户代码记录阶段AddNode 建立 DAG实例化 Init()拓扑排序/分流/建池/捕获 packetloop[多次重放]hipGraphCreate / 流捕获hipGraphInstantiatehipGraphLaunchRun() 批量派发预建 packethipGraphExecDestroyhipGraphDestroy对应的用户 API 与内部动作阶段用户 API内部动作创建/记录hipGraphCreate/hipStreamBeginCapture…EndCapture构造Graph通过AddNode建立节点与依赖实例化hipGraphInstantiate由Graph生成GraphExecInit()预计算一切执行hipGraphLaunch调用GraphExec::Run()重放销毁hipGraphExecDestroy/hipGraphDestroy释放实例与蓝图1.6 本章小结HIP Graph 解决的核心痛点是重复、固定拓扑负载下的 CPU launch 开销。核心思想是一次记录、多次重放本质是空间换时间——把昂贵的准备工作从每次执行挪到一次实例化。三大核心类分工明确GraphNode节点/零件、Graph蓝图/图纸、GraphExec可执行实例/生产线。两个关键继承设计GraphExec是一种GraphChildGraphNode同时是节点和可执行图支撑图嵌套。一次生命周期是创建 → 记录 → 实例化 →多次执行 → 销毁。下一章我们进入数据结构节点与图深入GraphNode的字段设计、节点派生类家族以及Graph如何用邻接表表示 DAG 并完成拓扑排序。
ROCm HIP Graph 概念与机制分析- 第 1 章: 概念与整体架构
发布时间:2026/7/17 19:55:45
接下来基于rocm-systems库源码文件系统剖析 HIP Graph 的概念、数据结构与执行机制。共分 9 章概念与整体架构数据结构节点与图数据结构辅助资源管理生命周期构建与实例化调度执行一经典路径与 Segment 分段调度执行二AQL Packet 捕获与批量派发调度执行三同步计划 SyncPlan高级特性子图、图内存与图更新实战与总结没有跑过达模型应用所以这部分不太硬本章目标讲清楚 HIP Graph 是什么、为什么需要它以及hip_graph_internal.hpp里三大核心类是如何分工协作的。读完本章你会对整个 Graph 子系统建立起一张地图为后续深入各个模块打好基础。1.1 从一个性能问题说起先看一段最普通的 HIP 代码——它是深度学习、科学计算里最常见的模式一轮迭代里连续提交一批小算子。for(intiter0;iter10000;iter){kernelAg,b,0,stream(...);// 提交 1kernelBg,b,0,stream(...);// 提交 2hipMemcpyAsync(...,stream);// 提交 3kernelCg,b,0,stream(...);// 提交 4}每一次kernel或hipMemcpyAsync调用运行时都要做一整套 CPU 端工作参数校验、查找 kernel、构造 launch 配置创建一个amd::Command对象计算依赖、挂上 event wait list组装 AQL packet、写入硬件队列当单个 kernel 很小几微秒时这套 CPU 端组装的开销可能和 kernel 本身的执行时间相当甚至更大。于是 GPU 频繁饿肚子——算得快但 CPU 喂不上。这就是所谓的CPU launch overhead / dispatch bound问题。关键观察是上面这个循环里每一轮的拓扑结构完全一样变的只是数据。既然结构不变为什么每一轮都要重新组装一遍1.2 Graph 的核心思想一次记录多次重放HIP Graph对标 CUDA Graph给出的答案是把做什么和怎么执行彻底分离重复记录阶段把一批操作及依赖构建成 DAG实例化阶段一次性预计算调度/资源/packet执行阶段直接重放可循环上万次记录阶段把一轮迭代里的所有操作kernel、memcpy、memset……及它们之间的依赖关系记录成一张有向无环图DAG。每个操作是一个节点依赖是边。实例化阶段把这张图编译成一个可执行实例在这一步把所有能提前算的东西全部算好——拓扑排序、流分配、kernel 参数池、信号池、甚至 GPU packet 本身。执行阶段重放时几乎没有 CPU 端组装工作直接把预先准备好的 packet 批量丢给硬件。一句话概括用空间预处理和缓存换时间重放性能。这是理解整个 Graph 子系统所有设计的总纲——后面你会看到大量在实例化时预计算、在执行时零成本的手法。1.3 三大核心类蓝图、节点、实例打开hip_graph_internal.hpp最顶上是一组前向声明classGraph;classGraphNode;classGraphExec;typedefGraphNode*Node;整个子系统就是围绕这三个类展开的它们的分工可以用图纸 → 生产线 → 零件来类比类角色类比关键职责GraphNode节点零件一道工序描述单个操作负责CreateCommand/EnqueueCommandsGraph图蓝图图纸组织节点与依赖负责增删节点、拓扑排序、克隆、内存池GraphExec可执行实例生产线按图纸建好的产线预分配 stream/信号/packet负责Run真正执行GraphNode —— 节点基类GraphNode是所有节点的抽象基类。它持有一个节点在图中的全部结构信息classGraphNode:publichipGraphNodeDOTAttribute{protected:hipGraphNodeType type_;// 节点类型std::vectoramd::Command*commands_;// 该节点对应的底层命令std::vectorNodeedges_;// 后继出边std::vectorNodedependencies_;// 前驱入边int32_tstream_id_-1;// 执行时分配到的流int32_tsegment_id_-1;// 分段调度中的段 IDstd::vectoruint8_t*gpuPackets_;// 预捕获的 GPU packet// ...};注意edges_和dependencies_这一对——它们把 DAG 的邻接关系直接存在了节点里。图的连接靠AddEdgeDep这类方法维护voidAddEdgeDep(constNodechildNode){AddEdge(childNode);// this 记录后继childNode-AddDependency(this);// child 记录前驱}每种具体操作对应一个派生类比如GraphKernelNodekernel 启动、GraphMemcpyNode拷贝、GraphMemsetNode填充、ChildGraphNode子图等等。它们通过重写虚函数CreateCommand/EnqueueCommands/SetParams/clone来定制各自行为——这是典型的模板方法 多态设计第 2 章会展开整个节点家族。Graph —— 可修改的蓝图Graph代表一张可以继续编辑的图。它管理所有顶点并提供构建、查询、调度的能力classGraph{std::vectorNodevertices_;// 图中所有节点std::vectorNodetopoOrder_;// 拓扑排序结果hip::MemoryPool*mem_pool_;// 图关联的内存池// ...public:voidAddNode(constNodenode);boolTopologicalOrder(std::vectorNodeTopoOrder);Graph*clone()const;// ...};Graph是活的你可以往里加节点、删节点、连边。它本身不负责执行——执行是GraphExec的事。GraphExec —— 可执行实例GraphExec通过多继承同时是一个Graph它 clone 了原图的内容和一个引用计数对象classGraphExec:publicamd::ReferenceCountedObject,publicGraph{std::unordered_mapint,std::vectorhip::Stream*parallel_streams_;// 预分配的流GraphKernelArgManager*kernArgManager_;// kernel 参数池GraphSignalManager*signalManager_;// HW 信号池std::unordered_mapint,SegmentBatchsegmentBatches_;// 预捕获的 packet 批SyncPlan sync_plan_;// 预计算的同步方案public:hipError_tInit();hipError_tRun(hip::Stream*stream);// ...};对比GraphGraphExec多出来的成员几乎全是预分配 / 预计算的缓存流池、参数池、信号池、packet 、同步计划。这正是 1.2 节空间换时间思想的直接体现——这些昂贵的准备工作只在Init()里做一次之后每次Run()都直接复用。1.4 类之间的关系全景把节点家族和三大核心类放在一起继承关系如下«abstract»GraphNodeCreateCommand()EnqueueCommands()clone()GraphAddNode()TopologicalOrder()clone()GraphExecInit()Run()ChildGraphNodeGraphKernelNodeGraphMemcpyNodeGraphMemsetNodeGraphHostNodeGraphMemAllocNode这里有两个值得注意的设计点GraphExec继承自Graph可执行实例是一种图——它拥有图的全部结构再叠加执行所需的资源。ChildGraphNode同时继承GraphNode和GraphExec一个子图节点既是父图里的一个节点本身又是一张可执行的图。这种双重身份让图可以嵌套是实现层级化图的关键第 8 章详解。1.5 一次完整生命周期鸟瞰最后把用户 API 和内部类串起来看一个 Graph 从生到死的完整流程GraphExecGraph用户代码GraphExecGraph用户代码记录阶段AddNode 建立 DAG实例化 Init()拓扑排序/分流/建池/捕获 packetloop[多次重放]hipGraphCreate / 流捕获hipGraphInstantiatehipGraphLaunchRun() 批量派发预建 packethipGraphExecDestroyhipGraphDestroy对应的用户 API 与内部动作阶段用户 API内部动作创建/记录hipGraphCreate/hipStreamBeginCapture…EndCapture构造Graph通过AddNode建立节点与依赖实例化hipGraphInstantiate由Graph生成GraphExecInit()预计算一切执行hipGraphLaunch调用GraphExec::Run()重放销毁hipGraphExecDestroy/hipGraphDestroy释放实例与蓝图1.6 本章小结HIP Graph 解决的核心痛点是重复、固定拓扑负载下的 CPU launch 开销。核心思想是一次记录、多次重放本质是空间换时间——把昂贵的准备工作从每次执行挪到一次实例化。三大核心类分工明确GraphNode节点/零件、Graph蓝图/图纸、GraphExec可执行实例/生产线。两个关键继承设计GraphExec是一种GraphChildGraphNode同时是节点和可执行图支撑图嵌套。一次生命周期是创建 → 记录 → 实例化 →多次执行 → 销毁。下一章我们进入数据结构节点与图深入GraphNode的字段设计、节点派生类家族以及Graph如何用邻接表表示 DAG 并完成拓扑排序。