本篇从 HIP 用户 API 入口出发，逐层追到 libhsakmt/KFD，重点回答一个技术问题——VMM 接口族和
hipMalloc*族到底差在哪里、为什么需要它。
1. 核心结论
hipMemCreate/hipMemAddressReserve/hipMemMap/hipMemSetAccess是 HIP 的虚拟内存管理（VMM）API，对标 CUDA 的cuMemCreate/cuMemAddressReserve/cuMemMap/cuMemSetAccess。
它和hipMalloc的本质差异：hipMalloc一步完成「分配物理内存 + 建 VA + 建页表 + 全设备可访问」；VMM 把这四件事拆成四个独立、可分别控制的 API：
hipMemAddressReserve → 只要一段虚拟地址（VA），不含物理内存 hipMemCreate → 只分配物理内存，返回句柄，不含 VA hipMemMap → 把「物理句柄」绑定到「VA 区间」 hipMemSetAccess → 逐设备授予 R/W 权限（此步才真正建 GPU 页表）整条链路的分层与ROCm HIP 内存分配接口到 libhsakmt 的调用路径分析一致，只是后端走的是 ROCr 的 VMM 接口而非 memory pool：
hipMemCreate/hipMemAddressReserve/hipMemMap/hipMemSetAccess->HIP Runtime API 层(hip_vm.cpp)->hip::GenericAllocation/amd::Memory(virtualbuffer 或 phys buffer)->roc::Device::virtual{Alloc,Map,Unmap,Free}/SetMemAccess/deviceVmemAlloc->Hsa::vmem_*(ROCclr 对 ROCr 符号的动态包装)->hsa_amd_vmem_address_reserve/vmem_handle_create/vmem_map/vmem_set_access->ROCr core::Runtime::VMemory*->hsaKmtAllocMemoryAlign(OnlyAddress)/hsaKmtMemoryVaMap/DRM ioctl->KFD2. 为什么需要 VMM：与 hipMalloc* 的差异（核心技术问题）
2.1 hipMalloc 的「一步到位」模型
回顾ROCm HIP 内存分配接口到 libhsakmt 的调用路径分析：hipMalloc → ihipMalloc → SvmBuffer::malloc → roc::Device::deviceLocalAlloc → hsa_amd_memory_pool_allocate。一次调用里，物理内存、VA、GPU 页表、（默认）本设备可访问性全部就绪。用户只拿到一个void*，无法单独控制其中任何一步。
这带来几个hipMalloc做不到的事：
| 需求 | hipMalloc 为何做不到 |
|---|---|
| 预留一大段连续 VA，之后再逐步填物理内存 | 分配即绑定，VA 不可预留 |
| 把不同来源的物理块（不同 GPU / 跨进程 dmabuf）拼进同一段连续 VA | 一个指针对应一块 pool 分配 |
| 先分配物理内存，稍后再决定映射到哪个 VA、给哪些设备访问 | 物理与 VA 绑死 |
| 精确控制「哪个设备以什么权限访问」 | 默认全设备可访问，粒度粗 |
| 动态扩容一块 buffer 而保持指针不变（在预留 VA 内追加映射） | 无法原地扩展 |
2.2 VMM 的「四步解耦」模型
| 步骤 | hipMalloc | VMM |
|---|---|---|
| ① 物理内存 | hipMalloc内隐式完成 | hipMemCreate→ 返回hipMemGenericAllocationHandle_t |
| ② 虚拟地址 | 隐式，由 runtime 定 | hipMemAddressReserve→ 返回void* |
| ③ 绑定/页表 | 隐式，全设备 | hipMemMap（簿记）+hipMemSetAccess（真正建页表） |
| ④ 访问授权 | 隐式，全设备可访问 | hipMemSetAccess逐{device, flags}授权 |
2.3 对象模型差异
ROCm HIP 内存分配接口到 libhsakmt 的调用路径分析 强调两个桥接：hip::Device→amd::Context和void*→amd::Memory。VMM 在void*→amd::Memory这一侧多出一层拆分：
hipMalloc: void* ptr ──► 一个 amd::Buffer（物理+VA 合一） VMM: hipMemCreate ──► phys amd::Memory（ROCCLR_MEM_PHYMEM）+ hip::GenericAllocation ← handle hipMemAddressReserve ──► parent virtual amd::Memory（只有 VA，无物理） ← ptr hipMemMap ──► 在 parent VA 上切出 sub-buffer，交叉链接到 phys mem- 物理侧：
hipMemCreate造的是ROCCLR_MEM_PHYMEMbuffer，其getUserData().hsa_handle保存 ROCr 返回的hsa_amd_vmem_alloc_handle_t；再包一层hip::GenericAllocation作为用户可见 handle。 - VA 侧：
hipMemAddressReserve造的是一个kParent=true的 virtual buffer，只登记 VA。 - 映射：
hipMemMap在 parent VA 里切出 sub-buffer，与 phys mem 交叉链接（phys_mem_obj）。这就是为什么hipMemSetAccess/hipMemUnmap里要沿 sub-buffer 链遍历。
3. HIP 层入口模式
所有 VMM API 同样是「HIP_INIT_API+ 参数检查 + 直接实现」，但注意：这些 API 大多是同步直达后端，不走 stream/command 队列（见hip_vm.cpp注释「Direct synchronous path」）。
4. 逐接口调用路径
4.1hipMemAddressReserve—— 只要 VA
hipMemAddressReserve(ptr,size,alignment,addr,flags)// 校验 size 是 virtualMemAllocGranularityMinimum_ 的整数倍、alignment 是 2 的幂->g_devices[0]->devices()[0]->virtualAlloc(addr,size,alignment)// 单次调用即为所有设备预留roc::Device::virtualAlloc -> Hsa::vmem_address_reserve(&vptr, size, req_addr, 0) // = hsa_amd_vmem_address_reserve -> CreateVirtualBuffer(context, vptr, size, kParent=true) // 建 parent virtual amd::Memory -> 返回 mem->getSvmPtr()底层：VMemoryAddressReserve → hsaKmtAllocMemoryAlign(OnlyAddress|FixedAddress)，只占 VA 不分配物理。
注意：VA 预留对所有设备一次完成（用
g_devices[0]即可），因为 VA 空间是进程级的。
4.2hipMemCreate—— 只分配物理内存
hipMemCreate(handle,size,prop,flags)// prop->type ∈ {Pinned, Uncached}；location ∈ {Device, Host}ihipFlags=ROCCLR_MEM_PHYMEM[|CL_MEM_SVM_ATOMICS|ROCCLR_MEM_HSA_UNCACHED]amdContext=(location==Host)?host_context:curDevContext->amd::SvmBuffer::malloc(*amdContext,ihipFlags,size,align)// 带 PHYMEM 标志->phys_mem_obj=getMemoryObject(dev,ptr)->phys_mem_obj->getUserData().data=newhip::GenericAllocation(...)->*handle=(hipMemGenericAllocationHandle_t)that GenericAllocationROCCLR_MEM_PHYMEM标志让 ROCm 后端走物理内存句柄分配而非普通 pool alloc：
roc::Memory::create（PHYMEM 路径） -> roc::Device::deviceVmemAlloc(size, flags) -> 选 pool（uncached ? gpu_ext_fine_grained : gpuvm_segment_） -> Hsa::vmem_handle_create(pool, size, MEMORY_TYPE_PINNED, 0, &hsa_vmem_handle) // = hsa_amd_vmem_handle_create -> 返回 hsa_vmem_handle.handle → 存入 userData().hsa_handle底层：VMemoryHandleCreate → region->Allocate + KfdDriver::CreateShareableHandle。此时没有 VA、没有页表。
对比
hipMalloc走的是deviceLocalAlloc → hsa_amd_memory_pool_allocate；VMM 走的是deviceVmemAlloc → hsa_amd_vmem_handle_create。这是两族接口在后端的第一处分岔。
4.3hipMemMap—— 绑定物理句柄到 VA（仅簿记）
hipMemMap(ptr,size,offset=0,handle,flags=0)ga=(hip::GenericAllocation*)handle// 校验 owner_dev_id 合法、size 对齐 granularityga->retain()->dev->virtualMap(ptr,size,&ga->asAmdMemory())// 同步直达，不走 streamroc::Device::virtualMap -> MapMemObjBookkeeping(phys, va, size) // 在 parent VA 上切 sub-buffer -> Hsa::vmem_map(sub_va, size, offset, hsa_handle, 0) // = hsa_amd_vmem_map -> FinalizeMapMemObjBookkeeping(...) // 交叉链接 sub_obj ↔ phys底层：VMemoryHandleMap只做簿记，无 ioctl——真正建页表在下一步。
4.4hipMemSetAccess—— 逐设备授权（真正建页表）
hipMemSetAccess(ptr,size,desc[],count)ValidateSubBufferCoverage(ptr,size)// size 必须正好覆盖整数个 sub-bufferforeach desc[i]{location,flags}:dev=g_devices[desc[i].location.id]->dev->devices()[0]->SetMemAccess(ptr,size,access_flags,locationType)roc::Device::SetMemAccess desc.permissions = access_flags // None/ReadOnly/ReadWrite desc.agent_handle = (kDevice) ? getBackendDevice() : getCpuAgent() -> Hsa::vmem_set_access(va, size, &desc, 1) // = hsa_amd_vmem_set_access底层：VMemorySetAccess → VMemorySetAccessPerHandle，GPU 走KfdDriver::Map → hsaKmtMemoryVaMap建 GPU 页表，CPU 走os::MapMemory。这一步之后 GPU 才能真正访问该 VA。
4.5hipMemUnmap/hipMemAddressFree/hipMemRelease—— 逆向拆解
hipMemUnmap(ptr,size) Pass1: 沿 sub-buffer 链校验 phys_mem_obj/ga、收集需同步的设备 Pass2: 对每个相关设备 SyncAllStreams() // 避免 unmap 与在飞行的访问竞争 Pass3: 逐 sub-buffer: sub_dev->virtualUnmap(sub_va,sub_size) -> Hsa::vmem_unmap (= hsa_amd_vmem_unmap) -> ga->release() hipMemAddressFree(ptr,size) -> g_devices[0]->devices()[0]->virtualFree(ptr) -> DestroyVirtualBuffer + Hsa::vmem_address_free (= hsa_amd_vmem_address_free) hipMemRelease(handle) -> ga->release() // 引用计数，GenericAllocation 归零时释放 phys（deviceVmemRelease → vmem_handle_release）4.6 共享与查询接口
| HIP API | roc backend | ROCr |
|---|---|---|
hipMemExportToShareableHandle | ExportShareableVMMHandle | vmem_export_shareable_handle（dmabuf）/vmem_export_fabric_handle |
hipMemImportFromShareableHandle | ImportShareableVMMHandle → ImportShareableHSAHandle | vmem_import_shareable_handle/vmem_import_fabric_handle |
hipMemRetainAllocationHandle | 查phys_mem_obj->userData().data | （纯 runtime 侧引用计数） |
hipMemGetAllocationGranularity | dev_info.virtualMemAllocGranularity* | — |
hipMemGetAllocationPropertiesFromHandle | GenericAllocation::GetProperties() | — |
hipMemGetAccess | GetMemAccess | vmem_get_access |
5. 完整调用链总览
6. 典型使用序列（对照 hipMalloc）
# hipMalloc 一步到位 ptr = hipMalloc(size); // 物理+VA+页表+全设备可访问 ... use ptr ... hipFree(ptr); # VMM 四步解耦 hipMemAddressReserve(&ptr, size, 0, 0, 0); // ① 预留 VA hipMemCreate(&h, size, &prop, 0); // ② 分配物理 hipMemMap(ptr, size, 0, h, 0); // ③ 绑定 hipMemSetAccess(ptr, size, &desc, 1); // ④ 授权（建页表） ... use ptr ... hipMemUnmap(ptr, size); hipMemRelease(h); hipMemAddressFree(ptr, size);7. 调试建议（分层断点）
HIP 层：
hipMemAddressReserve / hipMemCreate / hipMemMap / hipMemSetAccess / hipMemUnmapROCclr 层（rocdevice.cpp）：
roc::Device::virtualAlloc / virtualFree roc::Device::deviceVmemAlloc / deviceVmemRelease roc::Device::virtualMap / virtualUnmap roc::Device::SetMemAccess / GetMemAccess roc::Device::ExportShareableVMMHandle / ImportShareableHSAHandleROCr 层：
hsa_amd_vmem_address_reserve / vmem_handle_create / vmem_map / vmem_set_access core::Runtime::VMemoryAddressReserve / VMemoryHandleCreate / VMemoryHandleMap / VMemorySetAccesslibhsakmt 层：
hsaKmtAllocMemoryAlign / hsaKmtMemoryVaMap / hsaKmtMemoryVaUnmap / hsaKmtFreeMemory推荐第一轮断点（看「预留→创建→映射→授权」四步分岔）：
hipMemAddressReserve->virtualAlloc->hsa_amd_vmem_address_reserve hipMemCreate->deviceVmemAlloc->hsa_amd_vmem_handle_create hipMemMap->virtualMap->hsa_amd_vmem_map hipMemSetAccess->SetMemAccess->hsa_amd_vmem_set_access->hsaKmtMemoryVaMap8. 总结
- 和
hipMalloc*的根本差异：hipMalloc把「物理 + VA + 页表 + 访问」四件事一次做完、粒度粗、指针不可控；VMM 用hipMemCreate/hipMemAddressReserve/hipMemMap/hipMemSetAccess把它们拆成四个可独立控制的步骤，从而支持「预留大段连续 VA、按需拼贴不同物理块、逐设备精确授权、跨进程共享」等hipMalloc无法实现的场景。 - 对象模型上：VMM 在
void*→amd::Memory桥接里把「物理」（ROCCLR_MEM_PHYMEMbuffer +GenericAllocationhandle）和「VA」（parent virtual buffer + sub-buffer）拆开，hipMemMap负责把两者交叉链接。 - 后端分岔点：
hipMalloc走deviceLocalAlloc → hsa_amd_memory_pool_allocate；VMM 走deviceVmemAlloc → hsa_amd_vmem_handle_create以及virtual* → hsa_amd_vmem_{address_reserve,map,set_access}。再往下（ROCr → libhsakmt → KFD）的细节见下一篇[分析中…]。
