AMD rocr-libhsakmt分析系列1-1：多GPU系统中跨卡VRAM访问机制深度分析

原创于 2025-11-24 06:30:00 发布 · 502 阅读

5 ·

CC 4.0 BY-SA版权

文章标签：

#服务器 #linux #驱动开发 #ROCR #hsakmt #多卡互联

GPU软硬件系统设计专栏收录该内容

4 篇文章

订阅专栏

1. 场景概述

在8个GPU算力卡的系统中，GPU1需要访问GPU2-8的VRAM显存，驱动是如何实现的呢。本文基于ROCr-Runtime的libhsakmt库源码，深入分析这一跨GPU内存访问的完整实现机制。也是节点HsaIoLinkProperties信息的应用。

先给出实现流程概览：通过P2P拓扑发现 → 统一虚拟地址空间 → 页表映射 → 硬件转发，实现零拷贝的跨GPU内存访问。

2. 关键数据结构

2.1 GPU内存管理器 (gpu_mem_t)

每个GPU维护的P2P连接信息：

typedef struct {
    uint32_t gpu_id;                       // 当前GPU标识符
    uint32_t node_id;                      // NUMA节点ID
    
    // 核心P2P拓扑信息
    uint32_t usable_peer_id_num;           // 可访问的对等GPU数量
    uint32_t *usable_peer_id_array;        // 可访问的对等GPU ID数组
    
    manageable_aperture_t gpuvm_aperture;  // GPU虚拟地址空间管理
    int drm_render_fd;                     // DRM内核通信句柄
} gpu_mem_t;

2.2 内存对象结构 (vm_object_t)

每个显存分配的元数据：

struct vm_object {
    void *start;                           // 统一虚拟地址
    uint64_t size;                         // 分配大小
    uint32_t node_id;                      // 所属GPU的节点ID
    uint64_t handles[];                    // 内核驱动句柄
    
    // 映射状态追踪
    uint32_t *mapped_device_id_array;      // 已映射到哪些GPU
    uint32_t mapped_device_id_array_size;
    uint32_t mapping_count;                // 映射引用计数
};

3. 完整实现流程

3.1 阶段一：拓扑发现（系统初始化）

目标：确定哪些GPU之间可以互相访问。

// 在 hsakmt_fmm_init_process_apertures() 中执行
for (i = 0; i < num_sysfs_nodes; i++) {
    // 1. 获取当前GPU的IO链路信息
    ret = hsakmt_topology_get_node_props(nodeId, &nodeProps);
    ret = hsakmt_topology_get_iolink_props(nodeId, nodeProps.NumIOLinks, linkProps);
    
    // 2. 遍历所有IO链路，构建P2P连接表
    for (j = 0; j < nodeProps.NumIOLinks; j++) {
        int32_t to_gpu_mem_id = gpu_mem_find_by_node_id(linkProps[j].NodeTo);
        
        // 3. 关键：将当前GPU添加到对端GPU的可访问列表
        uint32_t peer = gpu_mem[to_gpu_mem_id].usable_peer_id_num++;
        gpu_mem[to_gpu_mem_id].usable_peer_id_array[peer] = gpu_mem[i].gpu_id;
    }
}

工作原理：

通过sysfs读取每个GPU的IO链路属性（/sys/class/kfd/kfd/topology/nodes/X/io_links/Y/）
若GPU A到GPU B存在IOLink，则在GPU B的usable_peer_id_array中添加GPU A的ID
最终每个GPU都知道"哪些GPU可以访问我的显存"

8卡系统示例（假设XGMI全连接）：

初始化后：
gpu_mem[0].usable_peer_id_array = [GPU1, GPU2, GPU3, GPU4, GPU5, GPU6, GPU7]
gpu_mem[1].usable_peer_id_array = [GPU0, GPU2, GPU3, GPU4, GPU5, GPU6, GPU7]
gpu_mem[2].usable_peer_id_array = [GPU0, GPU1, GPU3, GPU4, GPU5, GPU6, GPU7]
...

3.2 阶段二：显存分配（在GPU2上分配）

// 应用调用：在GPU2上分配1GB显存
void *gpu2_mem = hsakmt_fmm_allocate_device(
    gpu2_id, node2_id, NULL, 1GB, PAGE_SIZE, mflags);

// 内部实现流程：
void *hsakmt_fmm_allocate_device(...) {
    // 1. 选择GPU2的虚拟地址空间
    manageable_aperture_t *aperture = &gpu_mem[gpu2_mem_id].gpuvm_aperture;
    
    // 2. 分配虚拟地址（例如：0x7f8000000000）
    void *mem = aperture_allocate_area_aligned(aperture, NULL, size, alignment);
    
    // 3. 通过IOCTL在GPU2上分配物理显存
    vm_object_t *obj = fmm_allocate_memory_object(gpu2_id, mem, size, aperture, ...);
    
    // 4. 初始化映射状态为空
    obj->mapped_device_id_array = NULL;  // 尚未映射到任何GPU
    obj->mapping_count = 0;
    
    return mem;  // 返回虚拟地址 0x7f8000000000
}

关键点：

分配的是虚拟地址（所有GPU共享同一虚拟地址空间）
物理显存位于GPU2的VRAM
此时GPU1无法访问，因为GPU1的页表中还没有此地址的映射

3.3 阶段三：建立映射（GPU1访问GPU2显存）

用户态入口：

// 应用调用：将GPU2的显存映射到所有可访问的GPU
HSAKMT_STATUS ret = hsakmt_fmm_map_to_gpu(gpu2_mem, 1GB, NULL);

**IOCTL数据结构**：

```c
struct kfd_ioctl_map_memory_to_gpu_args {
    __u64 handle;                // GPU2显存的内核句柄
    __u64 device_ids_array_ptr;  // 指向目标GPU数组 [GPU0,GPU1,GPU3-7]
    __u32 n_devices;             // 要映射的GPU数量 = 7
    __u32 n_success;             // 内核返回：实际成功映射的GPU数量
};


**映射决策逻辑**：

```c
HSAKMT_STATUS hsakmt_fmm_map_to_gpu(void *address, uint64_t size, ...) {
    // 1. 查找内存对象（在GPU2的aperture中）
    vm_object_t *obj = vm_find_object(address, size, &aperture);
    
    // 2. 确定要映射到哪些GPU
    int32_t gpu_mem_id = gpu_mem_find_by_node_id(obj->node_id);  // 找到GPU2
    
    // 3. 使用GPU2的usable_peer_id_array作为映射目标
    uint32_t *target_gpus = gpu_mem[gpu_mem_id].usable_peer_id_array;
    uint32_t n_devices = gpu_mem[gpu_mem_id].usable_peer_id_num;
    // target_gpus 现在包含 [GPU0, GPU1, GPU3, GPU4, GPU5, GPU6, GPU7]
    
    // 4. 执行实际映射
    return _fmm_map_to_gpu(aperture, address, size, obj, target_gpus, n_devices);
}

核心映射实现：

static HSAKMT_STATUS _fmm_map_to_gpu(manageable_aperture_t *aperture,
                                     void *address, uint64_t size, vm_object_t *obj,
                                     uint32_t *nodes_to_map, uint32_t nodes_array_size)
{
    struct kfd_ioctl_map_memory_to_gpu_args args = {0};
    
    // 1. 准备IOCTL参数
    args.device_ids_array_ptr = (uint64_t)nodes_to_map;  
    // 指向 [GPU0, GPU1, GPU3, GPU4, GPU5, GPU6, GPU7]
    args.n_devices = nodes_array_size / sizeof(uint32_t);  // 7个GPU
    
    // 2. 对内存对象的每个句柄执行映射
    for (i = 0; i < obj->handle_num; i++) {
        args.handle = obj->handles[i];  // GPU2显存的内核句柄
        
        // 3. 系统调用：告诉内核"将这块显存映射到这7个GPU"
        ret = ioctl(kfd_fd, AMDKFD_IOC_MAP_MEMORY_TO_GPU, &args);
        
        if (ret < 0) {
            // 失败则回滚
            return HSAKMT_STATUS_ERROR;
        }
    }
    
    // 4. 记录映射状态
    obj->mapped_device_id_array = realloc(...);  // 更新为包含GPU0,1,3-7
    obj->mapped_device_id_array_size = args.n_success * sizeof(uint32_t);
    obj->mapping_count = 1;
    
    return HSAKMT_STATUS_SUCCESS;
}

3.4 阶段四：内核驱动处理

内核KFD驱动收到AMDKFD_IOC_MAP_MEMORY_TO_GPU后执行：

1. 验证P2P连接
   - 检查GPU0-7与GPU2之间是否存在物理IOLink
   
2. 为每个目标GPU建立页表映射
   对 GPU1 而言：
   - 在GPU1的页表中添加条目：虚拟地址0x7f8000000000 -> GPU2物理地址
   - 标记此页为"远程页"，访问时通过XGMI/PCIe
   
3. 配置硬件转发
   - 设置GPU1的内存控制器，将该地址范围的访问路由到GPU2
   - IOMMU配置（若启用）
   
4. 刷新TLB
   - 使新的页表映射立即生效

3.5 阶段五：访问执行

映射完成后，GPU1访问GPU2显存：

GPU1执行kernel:
  mov r0, [0x7f8000000000]  // 读取GPU2显存地址
  
硬件处理流程：
  1. GPU1查询本地页表
     -> 发现0x7f8000000000映射到GPU2的物理地址
     
  2. GPU1内存控制器
     -> 识别为远程访问，通过XGMI链路发起读请求
     
  3. XGMI/PCIe总线
     -> 将请求路由到GPU2
     
  4. GPU2内存控制器
     -> 从本地VRAM读取数据
     
  5. 返回路径
     -> 数据通过XGMI返回GPU1
     
  6. GPU1寄存器
     -> r0 = 读取的数据

关键特性：

整个过程无CPU参与（零拷贝）
使用统一虚拟地址（GPU1和GPU2看到相同地址）
硬件自动路由（由GPU内存控制器和XGMI处理）

4. 核心机制总结

4.1 数据流图

[系统启动]
    ↓
解析sysfs IOLink拓扑
    ↓
构建 usable_peer_id_array
    ├─ gpu_mem[0] → [GPU1, GPU2, ..., GPU7]
    ├─ gpu_mem[1] → [GPU0, GPU2, ..., GPU7]
    └─ gpu_mem[2] → [GPU0, GPU1, GPU3, ..., GPU7]
    
[GPU2分配显存]
    ↓
分配虚拟地址: 0x7f8000000000
    ↓
创建 vm_object (handle=123, node_id=2, mapped=[])

[应用调用 map_to_gpu]
    ↓
查询 gpu_mem[2].usable_peer_id_array → [GPU0,1,3-7]
    ↓
IOCTL: MAP_MEMORY_TO_GPU
    ├─ handle = 123
    ├─ devices = [GPU0, GPU1, GPU3, GPU4, GPU5, GPU6, GPU7]
    └─ n_devices = 7
    ↓
[内核KFD驱动]
    ├─ 在GPU0页表中建立映射
    ├─ 在GPU1页表中建立映射  ← GPU1现在可以访问
    ├─ 在GPU3页表中建立映射
    ├─ ...
    └─ 在GPU7页表中建立映射
    ↓
更新 vm_object.mapped_device_id_array = [GPU0,1,3-7]

[GPU1访问GPU2显存]
GPU1 kernel 访问 0x7f8000000000
    ↓
GPU1页表查询 → 映射到GPU2物理地址
    ↓
XGMI硬件转发 → GPU2 VRAM
    ↓
数据返回GPU1

4.2 三大关键机制

1. 统一虚拟地址空间（UVA）

所有GPU共享47位虚拟地址空间（0 - 128TB）
GPU2分配的地址在GPU1上仍然有效
简化编程模型：无需地址转换

2. 基于拓扑的访问控制

usable_peer_id_array白名单机制
只有IOLink连接的GPU才能互访
防止非法跨设备访问

3. 硬件辅助的页表映射

内核驱动负责页表管理
GPU硬件负责地址转换和数据转发
零CPU开销，全硬件加速

4.3 关键函数调用链

应用层
  hsakmt_fmm_map_to_gpu(address, size, NULL)
      ↓
  vm_find_object() - 查找内存对象
      ↓
  确定目标GPU列表 (usable_peer_id_array)
      ↓
  _fmm_map_to_gpu(aperture, address, size, obj, target_gpus, count)
      ↓
  准备IOCTL参数
      ↓
  ioctl(kfd_fd, AMDKFD_IOC_MAP_MEMORY_TO_GPU, &args)
      ↓
内核层
  KFD驱动: amdkfd_ioctl_map_memory_to_gpu()
      ↓
  验证P2P连接
      ↓
  为每个目标GPU调用 amdgpu_amdkfd_gpuvm_map_memory_to_gpu()
      ↓
  更新GPU页表 (amdgpu_vm_bo_map)
      ↓
  刷新TLB
      ↓
  返回成功状态

5. 实际应用示例

场景：8卡训练中的梯度同步

// GPU2分配梯度缓冲区
void *gradient_buffer = hsakmt_fmm_allocate_device(
    gpu2_id, node2_id, NULL, 128MB, PAGE_SIZE, mflags);
// 返回地址: 0x7f8000000000

// 映射到所有GPU
hsakmt_fmm_map_to_gpu(gradient_buffer, 128MB, NULL);
// GPU0,1,3-7 现在都可以通过 0x7f8000000000 访问此缓冲区

// GPU1执行kernel，读取GPU2的梯度
__global__ void AllReduce() {
    float *remote_grad = (float*)0x7f8000000000;  // GPU2的显存
    float local_grad = ...;
    float sum = local_grad + remote_grad[tid];   // 直接读取，零拷贝
    // XGMI硬件自动处理跨设备访问
}