ROCm rocr-runtime 库详细解析目录

文档系列简介 本系列文档旨在帮助开发者系统学习 AMD ROCm HSA Runtime（异构系统架构运行时），涵盖从基础概念到高级特性的完整技术栈。HSA Runtime 是 ROCm 软件栈的核心组件，提供了 GPU 设备管理、内存管理、任务调度、信号同步等关键功能。

主要模块: Agent、Queue、Signal、Memory、Loader 难度分级: 🟢 基础 | 🟡 中级 | 🔴 高级

📚 学习路线图

第一部分：HSA基础 (1-5章)
↓
第二部分：核心组件 (6-12章)
↓
第三部分：内存管理 (13-17章)
↓
第四部分：任务调度 (18-22章)
↓
第五部分：高级特性 (23-28章)

第一部分：HSA基础篇 🟢

建立异构计算和HSA架构的基础认知

01-什么是HSA与异构计算

• 1.1 异构计算的演进历史
• 1.2 HSA规范的诞生背景
• 1.3 HSA的核心设计理念
  • 统一内存空间
  • 用户态队列
  • 平台无关性
• 1.4 HSA与CUDA/OpenCL的对比
• 1.5 ROCm软件栈中的HSA定位

02-HSA-Runtime架构概览

• 2.1 Runtime层次结构
  • 公共API层 (inc/hsa.h)
  • 核心实现层 (core/runtime/)
  • 驱动抽象层 (core/driver/)
• 2.2 主要组件关系图
• 2.3 初始化与销毁流程
• 2.4 Runtime配置与环境变量
• 2.5 调试与追踪接口

03-HSA编程模型基础

• 3.1 Host与Device概念
• 3.2 执行模型
  • Kernel执行
  • Wavefront与Work-item
  • Work-group组织
• 3.3 内存模型
  • Global Memory
  • Group/Local Memory
  • Private Memory
• 3.4 同步原语
  • Signal
  • Barrier
  • Fence
• 3.5 HSA程序生命周期

04-开发环境搭建

• 4.1 ROCm安装与配置
• 4.2 HSA Runtime编译
• 4.3 示例程序运行
• 4.4 调试工具介绍
  • rocgdb
  • rocprof
  • HSA_DEBUG环境变量
• 4.5 常见问题排查

05-HSA-API快速入门

• 5.1 Runtime初始化hsa_init()
  hsa_shut_down()
• 5.2 设备发现与查询hsa_iterate_agents()
  hsa_agent_get_info()
• 5.3 内存分配hsa_memory_allocate()
  hsa_memory_free()
• 5.4 队列创建与管理hsa_queue_create()
  hsa_queue_destroy()
• 5.5 完整示例代码解析

第二部分：核心组件篇 🟡

深入理解Runtime的核心数据结构和机制

06-Agent设备抽象

• 6.1 Agent概念与类型
  • CPU Agent
  • GPU Agent
  • AIE Agent (AI Engine)
• 6.2 Agent类层次结构core::Agent (基类)
  ├── CpuAgent
  ├── GpuAgent/GpuAgentInt
  └── AieAgent
• 6.3 Agent属性查询
• 6.4 设备拓扑与互连
• 6.5 Agent生命周期管理

07-GPU-Agent深度解析

• 7.1 GpuAgent初始化流程
• 7.2 GPU硬件能力查询
  • Compute Units (CU)
  • Wavefront大小
  • ISA支持
• 7.3 GPU内存区域管理
• 7.4 DMA引擎 (SDMA/Blit)
• 7.5 GPU特有功能
  • Shader可见缓存
  • Trap Handler
  • Profiling支持

08-CPU-Agent与混合调度

• 8.1 CpuAgent实现特点
• 8.2 CPU内存访问
• 8.3 CPU-GPU协同
• 8.4 NUMA亲和性
• 8.5 Host Queue机制

09-Driver驱动抽象层

• 9.1 Driver接口设计
• 9.2 KFD (Kernel Fusion Driver) 集成
  • libhsakmt封装
  • Thunk层实现
• 9.3 设备发现与枚举
• 9.4 内存管理接口
• 9.5 队列管理接口

10-Topology拓扑发现

• 10.1 系统拓扑解析
• 10.2 NUMA节点映射
• 10.3 设备互连关系
  • PCIe链路
  • XGMI/Infinity Fabric
• 10.4 Link属性查询
• 10.5 拓扑信息缓存

11-ISA指令集架构

• 11.1 ISA抽象与版本管理
• 11.2 GCN/RDNA/CDNA架构差异
• 11.3 ISA能力查询
• 11.4 代码对象加载
• 11.5 向前兼容性处理

12-Runtime生命周期管理

• 12.1 初始化序列详解Runtime::Acquire() 流程
  ├── 加载Driver
  ├── 枚举Agents
  ├── 初始化MemoryRegions
  ├── 建立拓扑关系
  └── 注册Extension
• 12.2 延迟初始化 (Lazy Init)
• 12.3 多线程安全保证
• 12.4 资源清理与释放
• 12.5 异常处理机制

第三部分：内存管理篇 🟡🔴

掌握HSA的统一内存模型与分配策略

13-Memory-Region内存区域

• 13.1 MemoryRegion类设计
• 13.2 内存区域类型
  • System Memory (Host)
  • Local Memory (VRAM)
  • LDS (Local Data Share)
  • GDS (Global Data Share)
• 13.3 内存属性
  • Fine-grained vs Coarse-grained
  • Coherency特性
• 13.4 跨设备访问能力
• 13.5 内存池 (Pool) 管理

14-内存分配与释放

• 14.1 分配流程
• 14.2 内存对齐要求
• 14.3 大页支持
• 14.4 VRAM管理策略
• 14.5 内存碎片处理

15-统一内存与一致性

• 15.1 HSA统一地址空间
• 15.2 Fine-grained内存
  • 硬件缓存一致性
  • 原子操作支持
• 15.3 Coarse-grained内存
  • 显式同步需求
  • Fence与Acquire/Release
• 15.4 内存一致性模型
• 15.5 跨Agent内存访问

16-IPC共享内存

• 16.1 进程间内存共享

• 16.2 IPC Handle机制

• 16.3 GPU Direct RDMA

• 16.4 多进程场景

• 16.5 安全性考虑

17-DMA数据传输

• 17.1 DMA传输类型
  • Host ↔ Device
  • Device ↔ Device
  • Peer-to-Peer
• 17.2 SDMA引擎使用
• 17.3 Blit Kernel备用方案
• 17.4 异步拷贝与信号
• 17.5 性能优化技巧

文件:

• core/runtime/amd_blit_sdma.cpp
• core/runtime/amd_blit_kernel.cpp
• core/inc/blit_kernel.h

第四部分：任务调度篇 🟡🔴

理解AQL队列与内核调度机制

18-Queue队列机制

• 18.1 HSA队列模型
• 18.2 Queue类层次
• 18.3 用户态队列优势
• 18.4 队列创建参数
• 18.5 队列状态管理

19-AQL协议详解

• 19.1 AQL (Architected Queueing Language)

• 19.2 AQL数据包格式

• 19.3 Packet类型

  • Kernel Dispatch
  • Barrier
  • Agent Dispatch

• 19.4 Doorbell机制

• 19.5 Packet提交流程

20-Kernel-Dispatch内核派发

• 20.1 Dispatch Packet结构
• 20.2 工作项组织
  • Grid维度
  • Workgroup维度
  • Wavefront分配
• 20.3 内核参数传递
• 20.4 内存段配置
  • Private Scratch
  • Group LDS
• 20.5 完整派发示例

21-Scratch内存管理

• 21.1 Scratch的作用
  • 私有栈空间
  • Spill寄存器
• 21.2 动态Scratch分配
• 21.3 Scratch Cache优化
• 21.4 大Scratch处理
• 21.5 性能影响分析

22-Barrier与同步

• 22.1 Barrier Packet
  • AND语义
  • OR语义
• 22.2 依赖信号管理
• 22.3 多队列同步
• 22.4 内存Fence
  • Acquire Fence
  • Release Fence
• 22.5 同步开销优化

第五部分：高级特性篇 🔴

掌握信号、加载器、异常处理等高级功能

23-Signal信号系统

• 23.1 Signal设计理念
• 23.2 Signal类层次
• 23.3 SharedSignal结构
• 23.4 信号值操作
  • 原子读写
  • CAS (Compare-And-Swap)
  • 算术操作
• 23.5 信号等待策略

24-Signal等待机制

• 24.1 Wait条件判断
• 24.2 等待实现策略
  • Active Polling (忙等)
  • MWAITX指令
  • Event驱动 (中断)
• 24.3 超时处理
• 24.4 多信号等待
• 24.5 性能权衡

25-Loader代码加载器

• 25.1 HSA代码对象格式
  • ELF结构
  • .note段
  • .text段 (代码)
  • .rodata段 (常量)
• 25.2 代码对象加载流程
• 25.3 符号解析与重定位
• 25.4 多设备代码管理
• 25.5 动态加载与卸载

26-Trap-Handler异常处理

• 26.1 GPU异常类型
  • Memory Fault
  • Illegal Instruction
  • Trap指令
• 26.2 Trap Handler安装
• 26.3 异常信息收集
• 26.4 调试支持
• 26.5 Wave保存与恢复

27-Extension扩展机制

• 27.1 HSA Extension系统
• 27.2 AMD专有扩展
• 27.3 Image扩展
• 27.4 Finalizer扩展
• 27.5 Profile扩展
• 27.6 扩展注册与查询

28-性能分析与调试

• 28.1 Profiling支持
  • 时间戳查询
  • 性能计数器
• 28.2 API Trace
  • 调用拦截
  • 参数记录
• 28.3 调试环境变量
• 28.4 Core Dump分析
• 28.5 常见问题排查

学习路径建议

🎯 快速入门路径 (1-2周)

目标: 能够编写基本的HSA程序

🚀 系统掌握路径 (3-4周)

目标: 深入理解Runtime实现

🔍 问题驱动路径

目标: 针对性解决具体问题

• 性能优化: 07→17→21→28章
• 内存问题: 13→14→15→16章
• 调度问题: 18→19→20→22章
• 调试排查: 04→26→28章
• 扩展开发: 09→25→27章

代码阅读顺序建议

第一阶段：公共API理解

# 1. HSA公共头文件
inc/hsa.h # 核心API定义
inc/hsa_ext_amd.h # AMD扩展API

# 2. 基础数据结构
inc/amd_hsa_signal.h # Signal结构
inc/amd_hsa_queue.h # Queue结构

第二阶段：核心类设计

# 3. 核心抽象类
core/inc/agent.h # Agent基类
core/inc/queue.h # Queue基类
core/inc/signal.h # Signal基类
core/inc/memory_region.h # MemoryRegion基类

# 4. Runtime主控
core/inc/runtime.h # Runtime类定义

第三阶段：具体实现

# 5. Agent实现
core/runtime/amd_gpu_agent.cpp # GPU Agent
core/runtime/amd_cpu_agent.cpp # CPU Agent

# 6. Queue实现
core/runtime/amd_aql_queue.cpp # AQL队列
core/runtime/queue.cpp # Queue基础

# 7. Signal实现
core/runtime/signal.cpp # Signal实现
core/runtime/default_signal.cpp # 默认信号
core/runtime/interrupt_signal.cpp # 中断信号

# 8. Runtime初始化
core/runtime/runtime.cpp # Runtime实现

第四阶段：专题深入

# 9. 内存管理
core/runtime/amd_memory_region.cpp
core/runtime/amd_blit_sdma.cpp

# 10. 代码加载
loader/executable.cpp
core/runtime/amd_hsa_loader.cpp

# 11. 扩展功能
core/runtime/hsa_ext_amd.cpp

术语速查表

关键数据结构速查

Agent相关

// 基类
class core::Agent {
uint32_t node_id_;
DeviceType device_type_;
hsa_agent_t public_handle_;
// 虚函数：DmaCopy, DmaCopyAsync, …
};

// GPU实现
class GpuAgent : public Agent {
uint32_t node_id_;
HsaNodeProperties properties_;
std::vector<MemoryRegion*> regions_;
// SDMA, Blit, ISA, …
};

Queue相关

// AQL队列
class AqlQueue : public Queue {
amd_queue_t amd_queue_; // 硬件队列结构
uint64_t* ring_buf_; // 环形缓冲区
uint32_t ring_size_;
Signal* doorbell_signal_;
};

Signal相关

// 共享信号结构
struct SharedSignal {
amd_signal_t amd_signal; // 包含value, kind等
Signal* core_signal;
uint64_t timestamps[];
};

核心API速查

初始化

hsa_status_t hsa_init(void);
hsa_status_t hsa_shut_down(void);

Agent管理

hsa_status_t hsa_iterate_agents(
hsa_status_t (*callback)(hsa_agent_t, void*),
void* data);

hsa_status_t hsa_agent_get_info(
hsa_agent_t agent,
hsa_agent_info_t attribute,
void* value);

内存管理

hsa_status_t hsa_memory_allocate(
hsa_region_t region,
size_t size,
void** ptr);

hsa_status_t hsa_memory_free(void* ptr);

hsa_status_t hsa_memory_copy(
void* dst,
const void* src,
size_t size);

队列管理

hsa_status_t hsa_queue_create(
hsa_agent_t agent,
uint32_t size,
hsa_queue_type_t type,
void (*callback)(hsa_status_t, hsa_queue_t*, void*),
void* data,
uint32_t private_segment_size,
uint32_t group_segment_size,
hsa_queue_t** queue);

hsa_status_t hsa_queue_destroy(hsa_queue_t* queue);

uint64_t hsa_queue_add_write_index_relaxed(
const hsa_queue_t* queue,
uint64_t value);

信号管理

hsa_status_t hsa_signal_create(
hsa_signal_value_t initial_value,
uint32_t num_consumers,
const hsa_agent_t* consumers,
hsa_signal_t* signal);

void hsa_signal_destroy(hsa_signal_t signal);

hsa_signal_value_t hsa_signal_wait_scacquire(
hsa_signal_t signal,
hsa_signal_condition_t condition,
hsa_signal_value_t compare_value,
uint64_t timeout_hint,
hsa_wait_state_t wait_state_hint);

贡献指南

本文档系列持续完善中，欢迎贡献！

• 发现错误: 提交Issue说明问题
• 改进文档: Pull Request补充内容
• 提供示例: 特别欢迎实际代码案例
• 提出建议: 讨论学习路径优化

参考资源

官方文档

• HSA Specification
• ROCm Documentation
• HSA Runtime API Reference