HSA Runtime 架构概览

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

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

1. 🔧 接口统一：为不同硬件提供统一的 C++ 接口
2. 🎯 类型隔离：通过 DriverType 区分不同的驱动后端
3. 🔄 参数转换：将 HSA 类型转换为底层驱动所需的类型
4. 📊 生命周期：管理设备连接、资源分配与释放

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 交互机制、实战案例等详见后续章节。

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                                         │
│ ↓ ↓                                                       │
│ ┌─────────────┐ ┌─────────────┐                          │
│ │ 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 接入的关键点：

硬件操作通过 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_);

各层职责总结：

• 公共 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,...) \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");

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

🎯 本章总结

核心要点回顾

1. 三层架构设计
   • 公共 API 层：C 接口，符合 HSA 规范
   • 核心实现层：C++ 实现，面向对象设计
   • 驱动抽象层：Thunk + KFD，硬件访问
2. 主要组件关系
   • Runtime 管理所有 Agent 和 MemoryRegion
   • Agent 拥有 Queue 和内存资源
   • Queue 通过 Signal 实现同步
   • 清晰的生命周期管理
3. 初始化流程
   • 加载 Thunk 库 → 枚举设备 → 创建 Agent
   • 初始化内存区域 → 建立拓扑 → 加载扩展
   • 支持引用计数，多次初始化
4. 配置与调试
   • 丰富的环境变量控制行为
   • 工具拦截机制实现 API 追踪
   • 内置 Trace 点和 Profiling 支持

关键数据结构

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

思考题

1. 为什么 HSA Runtime 使用三层架构，而不是直接在 API 层操作硬件？
   • 可移植性：API 层与硬件解耦，同一套 API 支持不同厂商
   • 可维护性：业务逻辑 (核心层) 与硬件操作 (驱动层) 分离
   • 稳定性：API 层 ABI 稳定，内部实现可以优化而不影响应用
   • 安全性：驱动层在内核态，提供权限控制和资源隔离
2. Runtime 为什么使用引用计数而不是简单的初始化/销毁？
   • 多库场景：应用可能链接多个都依赖 HSA 的库（如 HIP、OpenCL）
   • 避免冲突：各库独立调用 init/shutdown 不会互相干扰
   • 资源共享：多个组件可以共用同一个 Runtime 实例
   • 线程安全：引用计数配合锁保证多线程安全
3. 工具拦截机制如何实现零开销的可选追踪？
   • 按需加载：只有设置 HSA_TOOLS_LIB 时才加载工具
   • 函数指针替换：拦截层只是简单的函数调用，无额外开销
   • 编译优化：Release 版本可以移除所有 Trace 宏
   • 惰性激活：工具库可以根据条件动态开关追踪
4. 为什么需要 Driver 类，而不是直接调用 Thunk 接口？
   • 面向对象封装：Driver 提供 C++ 接口，比 C 接口更易用和安全
   • 多驱动支持：通过虚函数支持 KFD、XDNA、Virtio 等不同后端
   • 资源管理：Driver 类管理设备文件描述符生命周期
   • 接口统一：上层组件（Runtime、Agent）无需关心底层驱动差异
   • 错误处理：统一的异常转换和错误码映射
   • 测试友好：可以 Mock Driver 类进行单元测试