📋 本章学习目标
完成本章学习后，你将能够：
- ✅ 理解 HSA Runtime 的三层架构设计
C++AI
HSA Runtime 架构概览
HSA Runtime 采用三层架构设计（公共 API 层、核心实现层、驱动抽象层），实现了接口与实现的分离。核心组件包括 Agent、Queue、Signal 和 Memory，通过 Driver 类统一封装底层硬件操作。 Runtime 的初始化与销毁流程、引用计数机制、配置环境变量以及调试追踪接口，为理解异构计算系统架构奠定基础。
HSA Runtime 采用三层架构设计（公共 API 层、核心实现层、驱动抽象层），实现了接口与实现的分离。核心组件包括 Agent、Queue、Signal 和 Memory，通过 Driver 类统一封装底层硬件操作。 Runtime 的初始化与销毁流程、引用计数机制、配置环境变量以及调试追踪接口，为理解异构计算系统架构奠定基础。
完成本章学习后，你将能够：
HSA Runtime 采用三层架构设计，实现了接口与实现的分离，提供了良好的可扩展性。
┌─────────────────────────────────────────────────┐
│ 应用程序 (User Application) │
│ (C/C++, Python, HIP, OpenCL 等) │
└─────────────────────────────────────────────────┘
↓ 调用
┌─────────────────────────────────────────────────┐
│ 第一层：公共 API 层 (Public API Layer) │
│ ┌───────────────────────────────────────────┐ │
│ │ inc/hsa.h - HSA 标准 API │ │
│ │ inc/hsa_ext_amd.h - AMD 扩展 API │ │
│ │ inc/hsa_ext_image.h - Image 扩展 │ │
│ │ inc/hsa_ext_finalize.h - Finalizer 扩展 │ │
│ └───────────────────────────────────────────┘ │
│ │
│ 特点： │
│ • C 语言接口，ABI 稳定 │
│ • 遵循 HSA 规范 │
│ • 跨平台、跨厂商 │
└─────────────────────────────────────────────────┘
↓ 实现
┌─────────────────────────────────────────────────┐
│ 第二层：核心实现层 (Core Runtime Layer) │
│ ┌───────────────────────────────────────────┐ │
│ │ core/runtime/runtime.cpp │ │
│ │ - Runtime 主控类 │ │
│ │ - Agent 管理 │ │
│ │ - 资源协调 │ │
│ └───────────────────────────────────────────┘ │
│ ┌───────────────────────────────────────────┐ │
│ │ 核心组件实现 (C++) │ │
│ │ ├─ core/runtime/amd_gpu_agent.cpp │ │
│ │ ├─ core/runtime/amd_cpu_agent.cpp │ │
│ │ ├─ core/runtime/amd_aql_queue.cpp │ │
│ │ ├─ core/runtime/signal.cpp │ │
│ │ ├─ core/runtime/amd_memory_region.cpp │ │
│ │ └─ core/runtime/amd_topology.cpp │ │
│ └───────────────────────────────────────────┘ │
│ ┌───────────────────────────────────────────┐ │
│ │ Driver 抽象 (core/inc/driver.h) │ │
│ │ - 面向对象的驱动接口封装 │ │
│ │ - Driver 基类定义 │ │
│ │ - 支持多种驱动类型 (KFD/XDNA/Virtio) │ │
│ │ - 统一的硬件操作接口 │ │
│ └───────────────────────────────────────────┘ │
│ │
│ 特点： │
│ • C++ 实现，面向对象设计 │
│ • 设备无关的抽象 │
│ • 业务逻辑核心 │
└─────────────────────────────────────────────────┘
↓ 调用
┌─────────────────────────────────────────────────┐
│ 第三层：驱动抽象层 (Driver Abstraction Layer) │
│ ┌───────────────────────────────────────────┐ │
│ │ Thunk Layer (libhsakmt.so) │ │
│ │ - KFD 用户态封装 │ │
│ │ - 系统调用转换 │ │
│ │ - 设备枚举 │ │
│ └───────────────────────────────────────────┘ │
│ ┌───────────────────────────────────────────┐ │
│ │ KFD (Kernel Fusion Driver) │ │
│ │ - amdgpu 内核模块 │ │
│ │ - 硬件资源管理 │ │
│ │ - 内存管理 │ │
│ │ - 队列调度 │ │
│ └───────────────────────────────────────────┘ │
│ │
│ 特点： │
│ • 内核态驱动 │
│ • 直接操作硬件 │
│ • 权限管理 │
└─────────────────────────────────────────────────┘
↓ 控制
┌─────────────────────────────────────────────────┐
│ GPU 硬件 (AMD GPU Hardware) │
│ GCN/RDNA/CDNA 架构 │
└─────────────────────────────────────────────────┘
这一层定义了符合 HSA 规范的 C 语言接口，是应用程序的直接入口：
// inc/hsa.h - 主要 API 分类
// 1. Runtime 生命周期
hsa_status_t hsa_init(void);
hsa_status_t hsa_shut_down(void);
// 2. 系统与 Agent 信息
hsa_status_t hsa_system_get_info(hsa_system_info_t attribute, void* value);
hsa_status_t hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent, void* data), void* data);
hsa_status_t hsa_agent_get_info(hsa_agent_t agent, hsa_agent_info_t attribute, void* value);
// 3. 内存管理
hsa_status_t hsa_memory_allocate(hsa_region_t region, size_t size, void** ptr);
hsa_status_t hsa_memory_free(void* ptr);
// 4. 队列操作
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);
// 5. 信号同步
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);
hsa_signal_value_t hsa_signal_wait_acquire(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);
API 设计原则：
回调遍历: 用于枚举资源
hsa_iterate_agents([](hsa_agent_t agent, void* data){
// 处理每个 agent
return HSA_STATUS_SUCCESS; // 继续遍历
}, user_data);
状态码返回: 统一的错误处理
typedef enum {
HSA_STATUS_SUCCESS = 0,
HSA_STATUS_ERROR_INVALID_ARGUMENT = 1,
HSA_STATUS_ERROR_OUT_OF_RESOURCES = 3,
// ... 更多状态码
} hsa_status_t;
Opaque Handle: 使用句柄（handle）隐藏内部实现
typedef struct hsa_agent_s { uint64_t handle; } hsa_agent_t;
typedef struct hsa_queue_s* hsa_queue_t;
这一层是 Runtime 的业务逻辑核心，使用 C++ 实现面向对象设计：
核心类体系：
// core/inc/runtime.h
namespace core {
class Runtime {
public:
// 单例模式
static Runtime* runtime_singleton_;
static bool Acquire();
static bool Release();
// Agent 管理
const std::vector<core::Agent*>& agents() const;
std::vector<Agent*>& gpu_agents() { return gpu_agents_; }
std::vector<Agent*>& cpu_agents() { return cpu_agents_; }
// 内存区域
const std::vector<core::MemoryRegion*>& regions() const;
// 驱动管理 - 支持多种驱动后端
std::vector<std::unique_ptr<Driver>>& AgentDrivers() { return agent_drivers_; }
Driver* GetDriverByType(DriverType drv_type);
// 系统拓扑
void RegisterAgent(Agent* agent);
private:
// 驱动接口数组 - 可同时支持多种类型的驱动
// 例如：KFD(用于 GPU)、XDNA(用于 AIE) 等
std::vector<std::unique_ptr<Driver>> agent_drivers_;
// 按设备类型分类的 Agent 列表
std::vector<Agent*> gpu_agents_; // GPU agents
std::vector<Agent*> cpu_agents_; // CPU agents
std::vector<Agent*> aie_agents_; // AIE agents (使用 XDNA 驱动)
// 所有 Agent 的统一索引（按 node_id）
std::map<uint32_t, Agent*> agents_by_node_;
};
// core/inc/agent.h
class Agent : public Checked<0x3D82E96E87AB41B1> {
public:
enum DeviceType {
kAmdCpuDevice,
kAmdGpuDevice,
kAmdDspDevice
};
virtual hsa_status_t QueueCreate(/* ... */) = 0;
virtual hsa_status_t DmaCopy(/* ... */) = 0;
uint32_t node_id() const { return node_id_; }
DeviceType device_type() const { return device_type_; }
protected:
uint32_t node_id_;
DeviceType device_type_;
Driver* driver_; // 每个 Agent 持有对应驱动的引用
};
} // namespace core
设计特点：
Driver 类的角色：
core::Driver 类是核心实现层对底层驱动的封装，它：
// core/inc/driver.h
namespace core {
enum class DriverType {
XDNA = 0,
KFD,
#ifdef HSAKMT_VIRTIO_ENABLED
KFD_VIRTIO,
#endif
NUM_DRIVER_TYPES
};
class Driver {
public:
// 初始化与关闭
virtual hsa_status_t Init() = 0;
virtual hsa_status_t ShutDown() = 0;
// 系统信息查询
virtual hsa_status_t GetSystemProperties(HsaSystemProperties& sys_props) const = 0;
virtual hsa_status_t GetNodeProperties(HsaNodeProperties& node_props, uint32_t node_id) const = 0;
// 内存操作
virtual hsa_status_t AllocateMemory(const MemoryRegion& mem_region, /*...*/) = 0;
virtual hsa_status_t FreeMemory(void* mem, size_t size) = 0;
// 队列操作
virtual hsa_status_t CreateQueue(uint32_t node_id, HSA_QUEUE_TYPE type, /*...*/) const = 0;
virtual hsa_status_t DestroyQueue(HSA_QUEUEID queue_id) const = 0;
// 时钟与配置
virtual hsa_status_t GetClockCounters(uint32_t node_id, HsaClockCounters* clock_counter) const = 0;
protected:
const DriverType kernel_driver_type_;
const std::string devnode_name_; // 如 "/dev/kfd"
int fd_ = -1; // 设备文件描述符
};
} // namespace core
设计模式：
在介绍驱动抽象层之前，我们需要理解Driver 在 Runtime 中的角色。
从上面的 Runtime 类定义可以看到，Runtime 维护了一个 agent_drivers_ 数组：
class Runtime {
private:
// 驱动接口数组 - 支持同时管理多种驱动类型
std::vector<std::unique_ptr<Driver>> agent_drivers_;
};
为什么是数组而不是单个 Driver？
现代异构系统中可能同时存在多种类型的计算设备：
DriverType::KFD)DriverType::XDNA)DriverType::KFD_VIRTIO)Runtime 需要同时管理这些不同类型的驱动，因此使用数组来存储。
Runtime 提供方法根据驱动类型获取对应的 Driver 实例：
// 根据驱动类型获取 Driver
Driver* Runtime::GetDriverByType(DriverType drv_type) {
auto is_driver_type = [drv_type](const std::unique_ptr<Driver>& d) {
return d->kernel_driver_type_ == drv_type;
};
auto driver = std::find_if(agent_drivers_.begin(), agent_drivers_.end(), is_driver_type);
return (driver != agent_drivers_.end()) ? driver->get() : nullptr;
}
使用场景：
现在让我们深入了解 Driver 的具体设计。
core::Driver)Driver 类提供了统一的硬件操作接口，所有具体驱动都继承此基类：
// core/inc/driver.h - Driver 抽象基类
namespace core {
enum class DriverType {
XDNA = 0, // 用于 AIE 设备
KFD, // 用于 AMD GPU (通过/dev/kfd)
KFD_VIRTIO, // 用于虚拟化 GPU
NUM_DRIVER_TYPES
};
class Driver {
public:
Driver(DriverType type, std::string devnode_name)
: kernel_driver_type_(type), devnode_name_(devnode_name) {}
// 生命周期管理
virtual hsa_status_t Init() = 0;
virtual hsa_status_t Open() = 0; // 打开设备节点
virtual hsa_status_t Close() = 0; // 关闭设备
// 系统拓扑查询
virtual hsa_status_t GetSystemProperties(...) const = 0;
virtual hsa_status_t GetNodeProperties(...) const = 0;
// 硬件资源管理
virtual hsa_status_t AllocateMemory(...) = 0;
virtual hsa_status_t CreateQueue(...) const = 0;
virtual hsa_status_t GetClockCounters(...) const = 0;
// ... 更多接口
// 驱动类型标识
DriverType driver_type() const { return kernel_driver_type_; }
protected:
const DriverType kernel_driver_type_; // 驱动类型
const std::string devnode_name_; // 设备节点路径
int fd_ = -1; // 设备文件描述符
};
} // namespace core
Driver 的核心职责：
DriverType 区分不同的驱动后端了解了 Driver 的接口设计后，让我们看看它如何在 Runtime 中被创建和使用：
// Runtime 初始化时创建 Driver 实例
bool Runtime::Acquire() {
if (ref_count_ == 0) {
// 1. 根据系统配置创建所需的 Driver
// 例如：检测到 GPU 则创建 KFD Driver，检测到 AIE 则创建 XDNA Driver
// 创建 KFD Driver 用于 GPU
if (HasGpuDevices()) {
auto kfd_driver = std::make_unique<KfdDriver>(DriverType::KFD, "/dev/kfd");
kfd_driver->Init();
kfd_driver->Open();
agent_drivers_.push_back(std::move(kfd_driver));
}
// 创建 XDNA Driver 用于 AIE（如果存在）
if (HasAieDevices()) {
auto xdna_driver = std::make_unique<XdnaDriver>(DriverType::XDNA, "/dev/accel/accel0");
xdna_driver->Init();
xdna_driver->Open();
agent_drivers_.push_back(std::move(xdna_driver));
}
// 2. 枚举设备并创建 Agent
// 每个 Agent 获取其对应类型的 Driver 引用
for (auto& driver : agent_drivers_) {
CreateAgentsForDriver(driver.get());
}
}
ref_count_++;
return true;
}
// 为特定 Driver 创建 Agent
void Runtime::CreateAgentsForDriver(Driver* driver) {
HsaSystemProperties sys_props;
driver->GetSystemProperties(sys_props);
for (uint32_t node = 0; node < sys_props.NumNodes; node++) {
HsaNodeProperties node_props;
driver->GetNodeProperties(node_props, node);
Agent* agent = nullptr;
if (IsGpuNode(node_props)) {
// GPU Agent - 传入 KFD Driver 引用
agent = new GpuAgent(node, node_props, driver);
gpu_agents_.push_back(agent);
} else if (IsAieNode(node_props)) {
// AIE Agent - 传入 XDNA Driver 引用
agent = new AieAgent(node, node_props, driver);
aie_agents_.push_back(agent);
}
if (agent) {
agents_by_node_[node] = agent;
}
}
}
Agent 如何使用 Driver：
// core/runtime/amd_gpu_agent.cpp
namespace amd {
class GpuAgent : public core::Agent {
private:
core::Driver* driver_; // 持有对应 Driver 的引用
public:
GpuAgent(uint32_t node_id, const HsaNodeProperties& props, core::Driver* driver)
: node_id_(node_id), driver_(driver) {
// Agent 不拥有 Driver，只持有引用
// Driver 的生命周期由 Runtime 管理
}
// 所有硬件操作都通过 driver_ 完成
hsa_status_t QueueCreate(uint32_t size, hsa_queue_t** queue) override {
void* ring_buffer;
HSAuint64 queue_id;
// 调用 Driver 接口，Driver 负责与底层通信
return driver_->CreateQueue(node_id_, size, &queue_id, &ring_buffer);
}
hsa_status_t AllocateMemory(size_t size, void** ptr) override {
return driver_->AllocateMemory(node_id_, size, HSA_MEM_HEAP_TYPE_FB_PUBLIC, ptr);
}
};
} // namespace amd
关键设计要点：
unique_ptr）💡 详细内容：Driver 的完整接口定义、KFD 交互机制、实战案例等详见后续章节。
libhsakmt（原来也叫做 Thunk）层是用户态的 C 接口库，封装了与 KFD 内核驱动的交互：
// 示例：Thunk 提供的接口
HSAKMT_STATUS HSAKMTAPI hsaKmtOpenKFD(void);
HSAKMT_STATUS HSAKMTAPI hsaKmtGetNodeProperties(HSAuint32 NodeId, HsaNodeProperties* NodeProperties);
HSAKMT_STATUS HSAKMTAPI hsaKmtAllocMemory(HSAuint32 PreferredNode, HSAuint64 SizeInBytes,
HsaMemFlags MemFlags, void** MemoryAddress);
HSAKMT_STATUS HSAKMTAPI hsaKmtCreateQueue(HSAuint32 NodeId, HSA_QUEUE_TYPE Type,
HSAuint32 QueuePercentage, HSA_QUEUE_PRIORITY Priority,
void* QueueAddress, HSAuint64 QueueSizeInBytes,
HsaEvent* Event, HsaQueueResource* QueueResource);
KFD (Kernel Fusion Driver): 内核态驱动
drivers/gpu/drm/amd/amdkfd/ioctl 与用户态通信三者关系：
Runtime/Agent (C++)
↓ 调用
core::Driver (C++ 抽象)
↓ 调用
Thunk Layer (libhsakmt.so, C 接口)
↓ ioctl 系统调用
KFD (内核驱动)
↓ 操作 GPU 硬件
┌────────────────────────────────────────────────────────────┐
│ Runtime (单例) │
│ ┌──────────────────────────────────────────────────────┐ │
│ │ • 全局状态管理 │ │
│ │ • Agent 分类管理 (GPU/CPU/AIE) │ │
│ │ • MemoryRegion 池 │ │
│ │ • Extension 管理 │ │
│ └──────────────────────────────────────────────────────┘ │
│ │
│ 持有（unique_ptr） │
│ ↓ │
│ ┌──────────────────────────────────────┐ │
│ │ agent_drivers_: vector<unique_ptr> │ │
│ │ ├─ KfdDriver (for GPU) │ │
│ │ ├─ XdnaDriver (for AIE) │ │
│ │ └─ VirtioDriver (for VM) │ │
│ └──────────────────────────────────────┘ │
│ │
│ 共享引用 | 管理 │
│ ↓ ↓ │
│ ┌─────────────┐ ┌─────────────┐ │
│ │ GPU Agents │ │ AIE Agents │ │
│ └─────────────┘ └─────────────┘ │
│ │
│ 通过各自的 Driver │
│ ↓ ↓ │
│ ┌─────────────┐ ┌─────────────┐ │
│ │ Thunk (KFD) │ │ Thunk (XDNA)│ │
│ └─────────────┘ └─────────────┘ │
└────────────────────────────────────────────────────────────┘
│ │
│ ioctl │
│ ioctl │
│ ↓ ↓ │
│ /dev/kfd /dev/accel/accel0 │
┌──────────────────────────────────────┐
│ Agent (抽象基类) │
│ ┌────────────────────────────────┐ │
│ │ 成员：Driver* driver_ │ ← 持有 Driver 引用 │
│ │ (不拥有，由 Runtime 管理) │ │
│ └────────────────────────────────┘ │
│ │
│ ├─ GpuAgent (使用 KfdDriver) │
│ │ - driver_->CreateQueue() │
│ │ - driver_->AllocateMemory() │
│ │ - VRAM/LDS 管理 │
│ │ │
│ ├─ CpuAgent (无需驱动) │
│ │ - System Memory │
│ │ - Host Queue │
│ │ │
│ └─ AieAgent (使用 XdnaDriver) │
│ - driver_->CreateQueue() │
│ - AIE 特定操作 │
└──────────────────────────────────────┘
│
│ 拥有 │
│ ↓ │
┌──────────────────────────────────────┐
│ MemoryRegion (内存区域) │
│ ├─ System Region (Fine/Coarse) │
│ ├─ Local Region (VRAM) │
│ ├─ LDS Region │
│ └─ Group Memory │
└──────────────────────────────────────┘
│
│ 分配 │
│ ↓ │
┌──────────────────────────────────────┐
│ Memory Allocation │
│ (用户的内存指针) │
└──────────────────────────────────────┘
┌──────────────────────────────────────┐
│ Queue (队列抽象) │
│ ├─ AqlQueue (GPU 硬件队列) │
│ ├─ HostQueue (CPU 软件队列) │
│ └─ InterceptQueue (调试/Trace) │
└──────────────────────────────────────┘
│
│ 包含 │
│ ↓ │
┌──────────────────────────────────────┐
│ AQL Packet Ring Buffer │
│ [Dispatch][Barrier][Dispatch]... │
└──────────────────────────────────────┘
│
│ 关联 │
│ ↓ │
┌──────────────────────────────────────┐
│ Signal (同步信号) │
│ ├─ DefaultSignal (共享内存) │
│ ├─ InterruptSignal (事件驱动) │
│ └─ IpcSignal (跨进程) │
└──────────────────────────────────────┘
应用程序启动
↓ hsa_init()
↓ Runtime::Acquire()
├─ 创建 Driver 实例
│ ├─ 根据平台选择 DriverType (KFD/XDNA)
│ └─ Driver::Open() 打开 /dev/kfd 或其他设备节点
│ ├─ 加载 Thunk 库 (通过 Driver)
├─ 枚举系统节点 (通过 Driver 接口)
│ ↓
│ 对每个节点创建 Agent
│ ├─ CPU Agent → System Memory Region
│ └─ GPU Agent → VRAM Region + LDS Region
│ (Agent 持有 Driver 引用，所有硬件操作通过 Driver)
│ ├─ 建立拓扑关系
│ (Agent 间的互连、NUMA 亲和性)
│ └─ 注册 Extensions (AMD 扩展、Image 扩展等)
应用程序运行
↓ hsa_queue_create()
↓ Agent::QueueCreate()
├─ 分配 Ring Buffer 内存 (通过 Driver::AllocateMemory)
├─ 向 KFD 注册队列 (通过 Driver::CreateQueue)
└─ 创建 Doorbell Signal
应用程序派发任务
↓ 写入 AQL Packet 到 Queue
↓ Doorbell 通知 GPU
↓ GPU 执行 Kernel
↓ Signal 同步等待
应用程序退出
↓ hsa_shut_down()
↓ Runtime::Release()
├─ 销毁所有 Queue (通过 Driver::DestroyQueue)
├─ 释放所有 Memory (通过 Driver::FreeMemory)
├─ 注销 Agents
├─ Driver::Close() 关闭设备
└─ 销毁 Driver 实例
以内存分配为例，展示各层的交互及 Driver 的接入：
应用程序调用
↓ hsa_memory_allocate(region, size, &ptr)
↓ [公共 API 层 - hsa.h]
验证参数，查找 MemoryRegion 对象
↓ [核心实现层 - MemoryRegion]
MemoryRegion::Allocate(size, &ptr)
↓ [核心实现层 - Agent]
Agent::AllocateMemory(size, &ptr)
├─ Agent 持有 driver_ 成员
└─ 通过 driver_ 调用底层
↓ [核心实现层 - Driver]
◄─── Driver 接入点
Driver::AllocateMemory(node_id, size, flags, &ptr)
├─ 参数转换：HSA 类型 → Thunk 类型
├─ 准备 HsaMemFlags 结构
└─ 调用 Thunk 接口
↓ [Thunk 层 - libhsakmt.so]
hsaKmtAllocMemory(NodeId, Size, Flags, &Address)
├─ 准备 ioctl 参数
└─ 执行系统调用
↓ ioctl(fd, AMDKFD_IOC_ALLOC_MEMORY_OF_GPU, &args)
↓ [KFD 内核驱动]
kfd_ioctl_alloc_memory_of_gpu()
├─ 权限检查
├─ 内存分配
└─ GPU 地址映射
↓ GPU 内存分配完成，返回地址
Driver 接入的关键点：
硬件操作通过 Driver：
// Agent 需要操作硬件时
hsa_status_t Agent::DoHardwareOperation() {
return driver_->SomeHardwareOp(node_id_, ...);
}
Agent 持有引用：
class Agent {
protected:
core::Driver* driver_; // 不拥有所有权，只持有引用
};
创建时接入：
// Runtime 初始化时创建 Driver
driver_ = CreateDriver(); // 根据平台选择 KfdDriver/XdnaDriver
// 创建 Agent 时传入 Driver 引用
agent = new GpuAgent(node_id, props, driver_);
各层职责总结：
// 应用程序调用
hsa_status_t status = hsa_init();
// ↓ 内部实现 (core/runtime/runtime.cpp)
bool Runtime::Acquire() {
ScopedAcquire<KernelMutex> lock(&runtime_lock_);
if (ref_count_ == 0) {
// 第一次初始化
// 1. 检测并创建所需的 Driver 实例
// 根据系统硬件配置，可能创建多个不同类型的 Driver
// 1a. 检测 GPU 设备，创建 KFD Driver
if (DetectGpuDevices()) {
auto kfd_driver = std::make_unique<KfdDriver>(DriverType::KFD, "/dev/kfd");
if (kfd_driver->Init() == HSA_STATUS_SUCCESS && kfd_driver->Open() == HSA_STATUS_SUCCESS) {
agent_drivers_.push_back(std::move(kfd_driver));
}
}
// 1b. 检测 AIE 设备，创建 XDNA Driver（如果存在）
if (DetectAieDevices()) {
auto xdna_driver = std::make_unique<XdnaDriver>(DriverType::XDNA, "/dev/accel/accel0");
if (xdna_driver->Init() == HSA_STATUS_SUCCESS && xdna_driver->Open() == HSA_STATUS_SUCCESS) {
agent_drivers_.push_back(std::move(xdna_driver));
}
}
// 2. 为每个 Driver 枚举并创建对应的 Agent
for (auto& driver : agent_drivers_) {
HsaSystemProperties sys_props;
driver->GetSystemProperties(sys_props);
// 枚举此驱动管理的所有节点
for (uint32_t node = 0; node < sys_props.NumNodes; node++) {
HsaNodeProperties node_props;
driver->GetNodeProperties(node_props, node);
Agent* agent = nullptr;
// 根据节点类型创建相应的 Agent
if (node_props.NumCPUCores > 0) {
// CPU 节点（不需要驱动）
agent = new amd::CpuAgent(node, node_props, nullptr);
cpu_agents_.push_back(agent);
} else if (node_props.NumFComputeCores > 0) {
// GPU 节点 - 传入 KFD Driver 引用
agent = new amd::GpuAgent(node, node_props, driver.get());
gpu_agents_.push_back(agent);
} else if (IsAieNode(node_props)) {
// AIE 节点 - 传入 XDNA Driver 引用
agent = new amd::AieAgent(node, node_props, driver.get());
aie_agents_.push_back(agent);
}
if (agent) {
// 将 Agent 添加到统一索引
agents_by_node_[node] = agent;
}
}
}
// 3. 为每个 Agent 初始化 MemoryRegion
// Agent 内部会通过 driver_ 查询内存属性
for (auto& [node_id, agent] : agents_by_node_) {
agent->InitRegions();
}
// 4. 建立 Agent 间的拓扑关系
DiscoverTopology();
// 5. 加载 Extensions
LoadExtensions();
// 6. 初始化工具库（Trace/Profile）
LoadTools();
}
ref_count_++;
return true;
}
初始化流程图：
hsa_init()
↓ Runtime::Acquire()
↓ ┌────────────────────────────────────┐
│ 1. 检测硬件并创建 Driver 实例 │
│ ┌──────────────────────────┐ │
│ │ 检测 GPU → KfdDriver │ │
│ │ → new KfdDriver("/dev/kfd") │ │
│ │ Init() + Open() │ │
│ │ agent_drivers_.push_back()│ │
│ └──────────────────────────┘ │
│ ┌──────────────────────────┐ │
│ │ 检测 AIE → XdnaDriver │ │
│ │ → new XdnaDriver("/dev/accel/accel0") │ │
│ │ Init() + Open() │ │
│ │ agent_drivers_.push_back()│ │
│ └──────────────────────────┘ │
└────────────────────────────────────┘
↓ ┌────────────────────────────────────┐
│ 2. 为每个 Driver 枚举设备节点 │
│ 对 agent_drivers_ 中每个 driver │
│ ↓ │
│ driver->GetSystemProperties() │
│ driver->GetNodeProperties() │
└────────────────────────────────────┘
↓ ┌────────────────────────────────────┐
│ 3. 创建 Agent 并关联 Driver │
│ ┌─────────────────────────────┐ │
│ │ CpuAgent(node, props, nullptr) │ → cpu_agents_ │
│ │ （CPU 不需要驱动） │ │
│ └─────────────────────────────┘ │
│ ┌─────────────────────────────┐ │
│ │ GpuAgent(node, props, kfd_driver) │ → gpu_agents_ │
│ │ （传入 KFD Driver 引用） │ │
│ └─────────────────────────────┘ │
│ ┌─────────────────────────────┐ │
│ │ AieAgent(node, props, xdna_driver) │ → aie_agents_ │
│ │ （传入 XDNA Driver 引用） │ │
│ └─────────────────────────────┘ │
│ 所有 Agent 统一索引到 │
│ agents_by_node_[node] │
└────────────────────────────────────┘
↓ ┌────────────────────────────────────┐
│ 4. 初始化 MemoryRegion │
│ agent->InitRegions() │
│ (Agent 通过 driver_ 查询内存) │
└────────────────────────────────────┘
↓ ┌────────────────────────────────────┐
│ 5. 建立拓扑关系 │
│ - Agent 间连接 │
│ - NUMA 亲和性 │
└────────────────────────────────────┘
↓ ┌────────────────────────────────────┐
│ 6. 加载 Extensions │
│ - AMD 扩展 │
│ - Image 扩展 │
└────────────────────────────────────┘
↓ ┌────────────────────────────────────┐
│ 7. 初始化工具 │
│ - API Trace │
│ - Profiler │
└────────────────────────────────────┘
↓ 返回 HSA_STATUS_SUCCESS
关键设计：
• agent_drivers_ 是 vector，支持多个 Driver 共存
• 每个 Agent 持有对应类型 Driver 的引用
• Driver 由 Runtime 统一管理生命周期
• 支持异构系统（GPU+AIE 同时存在）
// 应用程序调用
hsa_shut_down();
// ↓ 内部实现
bool Runtime::Release() {
ScopedAcquire<KernelMutex> lock(&runtime_lock_);
if (ref_count_ == 1) {
// 最后一次引用
// 1. 卸载工具库
UnloadTools();
// 2. 卸载 Extensions
UnloadExtensions();
// 3. 销毁所有 Agent（自动清理 Queue、Memory）
for (auto agent : agents_) {
delete agent;
}
agents_.clear();
gpu_agents_.clear();
cpu_agents_.clear();
// 4. 关闭 KFD 设备
hsaKmtCloseKFD();
// 5. 卸载 Thunk 库
Unload();
}
ref_count_--;
return true;
}
HSA Runtime 使用引用计数支持多次初始化/销毁：
// 场景：多个库都依赖 HSA Runtime
// 库 A 初始化
hsa_init(); // ref_count: 0→1, 实际初始化
// 库 B 初始化
hsa_init(); // ref_count: 1→2, 无操作
// 库 A 清理
hsa_shut_down(); // ref_count: 2→1, 无操作
// 库 B 清理
hsa_shut_down(); // ref_count: 1→0, 实际清理
HSA Runtime 提供丰富的环境变量来控制行为和调试。
| 环境变量 | 作用 | 取值 | 示例 |
|---|---|---|---|
HSA_ENABLE_DEBUG | 启用调试模式 | 0/1 | export HSA_ENABLE_DEBUG=1 |
HSA_TOOLS_LIB | 加载工具库 | 库路径 | export HSA_TOOLS_LIB=libhsa-runtime-tools64.so |
HSA_TOOLS_REPORT_LOAD_FAILURE | 报告工具加载失败 | 0/1 | export HSA_TOOLS_REPORT_LOAD_FAILURE=1 |
HSA_QUEUE_SIZE | 默认队列大小 | 数值 | export HSA_QUEUE_SIZE=4096 |
HSA_SCRATCH_MEM | Scratch 内存大小 (MB) | 数值 | export HSA_SCRATCH_MEM=256 |
HSA_RUNNING_KERNEL_MODE | 内核运行模式 | 0/1 | 用于调试 |
HSA_SVM_GUARD_PAGES | SVM 保护页 | 0/1 | 内存调试 |
HSA_CHECK_FLAT_SCRATCH | 检查 Scratch 访问 | 0/1 | 性能分析 |
# 1. API 追踪
export HSA_TOOLS_LIB=/opt/rocm/lib/libhsa-runtime-tools64.so # 输出所有 HSA API 调用和参数
# 2. 队列调试
export HSA_QUEUE_DEBUG=1 # 打印队列创建/销毁信息
# 3. 内存调试
export HSA_MEMORY_DEBUG=1 # 追踪内存分配/释放
# 4. Signal 调试
export HSA_SIGNAL_DEBUG=1 # 打印 Signal 等待/触发
# 5. 详细日志
export HSA_ENABLE_DEBUG=1
export HSA_DEBUG_LEVEL=7 # 0-7，数字越大越详细
# 1. Scratch 内存预分配
export HSA_SCRATCH_MEM=512 # 预分配 512MB Scratch，避免运行时分配
# 2. 队列大小
export HSA_QUEUE_SIZE=8192 # 更大的队列可以容纳更多任务
# 3. 禁用一致性检查（生产环境）
export HSA_CHECK_FLAT_SCRATCH=0
export HSA_ENABLE_DEBUG=0
# 4. SDMA 引擎优化
export HSA_ENABLE_SDMA=1 # 优先使用硬件 DMA 传输
部分配置可以通过配置文件 /etc/hsa/hsa.conf 设置：
# /etc/hsa/hsa.conf 示例
[runtime]
# 默认队列大小
queue_size = 4096
# Scratch 内存
scratch_mem_mb = 256
# 是否启用 SDMA
enable_sdma = true
[debug]
# 调试级别 0-7
level = 0
# API 追踪
api_trace = false
[memory]
# 大页支持
use_huge_pages = true
# SVM 保护页
svm_guard_pages = false
HSA Runtime 提供了工具拦截机制，允许外部库拦截所有 API 调用：
// HSA 定义的工具接口
typedef struct hsa_api_table_s {
hsa_status_t (*hsa_init_fn)(void);
hsa_status_t (*hsa_shut_down_fn)(void);
hsa_status_t (*hsa_agent_get_info_fn)(hsa_agent_t, hsa_agent_info_t, void*);
// ... 所有 HSA API 的函数指针
} hsa_api_table_t;
// 工具库导出的函数
extern "C" bool OnLoad(void* table, // Runtime 的原始函数表
uint64_t runtime_version,
uint64_t failed_tool_count,
const char* const* failed_tool_names) {
// 保存原始函数
original_table = *(hsa_api_table_t*)table;
// 替换为拦截函数
hsa_api_table_t* api_table = (hsa_api_table_t*)table;
api_table->hsa_init_fn = InterceptHsaInit;
api_table->hsa_agent_get_info_fn = InterceptHsaAgentGetInfo;
// ...
return true;
}
// 拦截函数示例
hsa_status_t InterceptHsaInit(void) {
printf("[Trace] hsa_init() called\n");
auto start = std::chrono::high_resolution_clock::now();
hsa_status_t status = original_table.hsa_init_fn();
auto end = std::chrono::high_resolution_clock::now();
printf("[Trace] hsa_init() returned %d, took %lld us\n", status, std::chrono::duration_cast<std::chrono::microseconds>(end - start).count());
return status;
}
使用工具库：
# 加载追踪工具
export HSA_TOOLS_LIB=/opt/rocm/lib/libhsa-runtime-tools64.so
# 运行应用
./my_hsa_app
# 输出示例：
# [Trace] hsa_init() called
# [Trace] hsa_init() returned 0, took 1234 us
# [Trace] hsa_iterate_agents() called
# [Trace] Agent 0: CPU, 16 cores
# [Trace] Agent 1: GPU, 60 CUs
# ...
Runtime 内部使用宏定义追踪点：
// core/inc/runtime.h
#ifdef HSA_TRACE
#define TRACE_RUNTIME(fmt,...) \n fprintf(stderr, "[HSA] " fmt "\n", ##__VA_ARGS__)
#else
#define TRACE_RUNTIME(...)
#endif
// 使用示例
TRACE_RUNTIME("Agent %u initialized: %s", agent->node_id(),
agent->device_type() == Agent::kAmdGpuDevice ? "GPU" : "CPU");
// HSA 提供的时间戳查询
hsa_status_t hsa_system_get_info(
HSA_SYSTEM_INFO_TIMESTAMP, // 获取硬件时间戳
×tamp);
// 高精度时间戳频率
hsa_system_get_info(
HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY,
&frequency);
// 计算时间
double seconds = (end_timestamp - start_timestamp) / (double)frequency;
// AMD 扩展：设置异常策略
hsa_amd_signal_async_handler(signal, HSA_SIGNAL_CONDITION_NE, 0, exception_handler, NULL);
// AMD 扩展：内存追踪
hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SIZE, &total_size);
hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_ALLOC_MAX_SIZE, &max_alloc);
// Runtime 核心类
class Runtime {
static Runtime* runtime_singleton_;
std::vector<Agent*> agents_;
std::vector<MemoryRegion*> regions_;
core::Driver* driver_; // 驱动实例
int ref_count_;
};
// Driver 抽象（位于核心实现层）
class Driver {
DriverType kernel_driver_type_; // KFD/XDNA/VIRTIO
std::string devnode_name_; // 设备节点路径
int fd_; // 设备文件描述符
// 提供所有硬件操作接口
};
// Agent 抽象
class Agent {
uint32_t node_id_;
DeviceType device_type_;
std::vector<MemoryRegion*> regions_;
core::Driver* driver_; // 持有 Driver 引用
};
HSA_TOOLS_LIB 时才加载工具
微信公众号「极客日志」，在微信中扫描左侧二维码关注。展示文案：极客日志 zeeklog
生成新的随机RSA私钥和公钥pem证书。 在线工具，RSA密钥对生成器在线工具，online
基于 Mermaid.js 实时预览流程图、时序图等图表，支持源码编辑与即时渲染。 在线工具，Mermaid 预览与可视化编辑在线工具，online
将字符串编码和解码为其 Base64 格式表示形式即可。 在线工具，Base64 字符串编码/解码在线工具，online
将字符串、文件或图像转换为其 Base64 表示形式。 在线工具，Base64 文件转换器在线工具，online
将 Markdown（GFM）转为 HTML 片段，浏览器内 marked 解析；与 HTML转Markdown 互为补充。 在线工具，Markdown转HTML在线工具，online
将 HTML 片段转为 GitHub Flavored Markdown，支持标题、列表、链接、代码块与表格等；浏览器内处理，可链接预填。 在线工具，HTML转Markdown在线工具，online