第02章：HSA-Runtime架构概览

章节概述 本章深入HSA Runtime的内部架构，介绍其三层结构设计、核心组件关系以及初始化流程。通过理解Runtime的整体架构，你将为后续学习具体模块打下坚实基础。

难度级别: 🟢 基础 预计阅读时间: 45分钟 前置知识: 第01章 – 什么是HSA与异构计算

📋 本章学习目标

完成本章学习后，你将能够：

• ✅ 理解HSA Runtime的三层架构设计
• ✅ 掌握主要组件（Agent、Queue、Signal、Memory）的关系
• ✅ 了解Runtime的初始化与销毁流程
• ✅ 熟悉Runtime配置选项和环境变量
• ✅ 知道如何启用调试和追踪功能

2.1 Runtime层次结构

HSA Runtime采用三层架构设计，实现了接口与实现的分离，提供了良好的可扩展性。

2.1.1 三层架构总览

┌─────────────────────────────────────────────────┐
│ 应用程序 (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架构 │
└─────────────────────────────────────────────────┘

2.1.2 公共API层 (inc/hsa.h)

这一层定义了符合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_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);

API设计原则：

• Opaque Handle: 使用句柄（handle）隐藏内部实现typedef struct hsa_agent_s { uint64_t handle; } hsa_agent_t;
  typedef struct hsa_queue_s* hsa_queue_t;
• 状态码返回: 统一的错误处理typedef enum {
  HSA_STATUS_SUCCESS = 0,
  HSA_STATUS_ERROR_INVALID_ARGUMENT = 1,
  HSA_STATUS_ERROR_OUT_OF_RESOURCES = 3,
  // … 更多状态码
  } hsa_status_t;
• 回调遍历: 用于枚举资源hsa_iterate_agents([](hsa_agent_t agent, void* data) {
  // 处理每个agent
  return HSA_STATUS_SUCCESS; // 继续遍历
  }, user_data);

2.1.3 核心实现层 (core/runtime/)

这一层是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

设计特点：

• Runtime使用驱动数组而非单一驱动，支持异构系统中的多种设备
• 每个Driver实例负责一类设备（如KFD负责GPU，XDNA负责AIE）
• Agent创建时从Runtime获取对应类型的Driver引用

Driver类的角色： core::Driver 类是核心实现层对底层驱动的封装，它：

• 提供统一的C++接口给上层组件（Runtime、Agent等）
• 支持多种驱动后端（KFD、XDNA、KFD_VIRTIO）
• 封装了所有与内核驱动交互的细节

// 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

设计模式：

• 工厂模式: Runtime根据硬件类型创建Agent和Driver
• 单例模式: Runtime全局唯一实例
• 策略模式: 不同类型的Queue、Signal、Driver实现
• 观察者模式: Signal等待机制

2.1.4 Driver的角色与定位

在介绍驱动抽象层之前，我们需要理解Driver在Runtime中的角色。

Driver在Runtime中的组织方式

从上面的Runtime类定义可以看到，Runtime维护了一个 agent_drivers_ 数组：

class Runtime {
private:
// 驱动接口数组 – 支持同时管理多种驱动类型
std::vector<std::unique_ptr<Driver>> agent_drivers_;
};

为什么是数组而不是单个Driver？

现代异构系统中可能同时存在多种类型的计算设备：

• GPU设备 → 使用KFD驱动 (DriverType::KFD)
• AIE设备 → 使用XDNA驱动 (DriverType::XDNA)
• 虚拟化环境 → 使用Virtio驱动 (DriverType::KFD_VIRTIO)

Runtime需要同时管理这些不同类型的驱动，因此使用数组来存储。

Driver的获取与使用

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;
}

使用场景：

• Agent创建时获取对应类型的Driver
• Runtime初始化时根据系统配置创建相应Driver
• 支持多种驱动共存于同一系统

2.1.5 驱动抽象层详解 (Driver + libhsakmt + KFD)

现在让我们深入了解Driver的具体设计。

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的核心职责：

Driver的初始化与Agent的关联

了解了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

关键设计要点：

• ✅ Runtime持有Driver的所有权（unique_ptr）
• ✅ Agent只持有Driver的引用（裸指针）
• ✅ 多个Agent可以共享同一个Driver实例
• ✅ Driver的生命周期由Runtime统一管理

💡 详细内容：Driver的完整接口定义、KFD交互机制、实战案例等详见 第09章 – Driver驱动抽象层

libhsakmt层与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): 内核态驱动

• 位于Linux内核的 drivers/gpu/drm/amd/amdkfd/
• 通过ioctl与用户态通信
• 管理GPU资源、内存、队列

三者关系：

Runtime/Agent (C++)
↓ 调用
core::Driver (C++ 抽象)
↓ 调用
Thunk Layer (libhsakmt.so, C接口)
↓ ioctl系统调用
KFD (内核驱动)
↓ 操作
GPU硬件

2.2 主要组件关系图

2.2.1 核心组件交互

┌────────────────────────────────────────────────────────────┐
│ 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 │通过各自的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 (跨进程) │
└──────────────────────────────────────┘

2.2.2 对象生命周期

应用程序启动
↓
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实例

2.2.3 典型调用链示例

以内存分配为例，展示各层的交互及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接入的关键点：

各层职责总结：

• 公共API层：标准接口，参数校验
• 核心实现层：业务逻辑，资源管理
• Driver层：接口适配，参数转换，硬件抽象
• Thunk层：系统调用封装
• KFD驱动：硬件操作，权限控制

2.3 初始化与销毁流程

2.3.1 初始化详细流程

// 应用程序调用
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同时存在）

2.3.2 销毁流程

// 应用程序调用
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;
}

2.3.3 引用计数机制

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, 实际清理

2.4 Runtime配置与环境变量

HSA Runtime提供丰富的环境变量来控制行为和调试。

2.4.1 常用环境变量

2.4.2 调试相关变量

# 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，数字越大越详细

2.4.3 性能优化变量

# 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传输

2.4.4 配置文件

部分配置可以通过配置文件 /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

2.5 调试与追踪接口

2.5.1 API追踪工具

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
# …

2.5.2 内部Trace点

Runtime内部使用宏定义追踪点：

// core/inc/runtime.h
#ifdef HSA_TRACE
#define TRACE_RUNTIME(fmt, ...) \\
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");

2.5.3 Profiling接口

// HSA提供的时间戳查询
hsa_status_t hsa_system_get_info(
HSA_SYSTEM_INFO_TIMESTAMP, // 获取硬件时间戳
&timestamp);

// 高精度时间戳频率
hsa_system_get_info(
HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY,
&frequency);

// 计算时间
double seconds = (end_timestamp – start_timestamp) / (double)frequency;

2.5.4 AMD扩展调试功能

// 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引用
};

思考题

🔗 导航

• 上一章: 01-什么是HSA与异构计算
• 下一章: 第03章：HSA编程模型基础
• 返回目录: 00-HSA-Runtime学习文档目录

📚 延伸阅读

ROCm rocr-libhsakmt分析系列0：HSA(异构系统架构)驱动概览

相关主题

• 设计模式：《Design Patterns》(GoF)
• 系统架构：《Software Architecture Patterns》
• 插件系统：《Plugin Architecture》
• 驱动开发：Linux Device Drivers (3rd Edition)