1. 核心结论HIP 内存分配链路的核心是看清楚用户态指针 API 如何逐层转换成 ROCclr memory object、ROCr HSA memory/SVM 操作以及最终的 libhsakmt/KFD 资源管理动作。主路径可以概括为hipMalloc/hipMallocManaged/hipHostMalloc-HIP Runtime API 层-ihipMalloc/ihipMallocManaged/ihipHostMalloc-hip::Device/amd::Context-amd::SvmBuffer/amd::Buffer/amd::Memory-roc::Memory/roc::Device-Hsa::memory_pool_allocate/Hsa::svm_attributes_set-ROCr Runtime-hsaKmtAllocMemory/hsaKmtMapMemoryToGPUNodes/hsaKmtSVMSetAttr-KFD ioctl这条链路里有三个关键转换阶段看到的对象说明HIP APIvoid*用户只看到裸指针HIP/ROCclramd::Memory/amd::Bufferruntime 用内存对象保存大小、flags、context、device 等元数据ROCr/libhsakmtHSA memory pool / SVM range / KFD allocation真正和 KFD 交互的底层资源根据当前 hip::Device 找到 amd::Context - 创建 ROCclr memory object - 在 ROCm 后端创建 roc::Memory - 通过 ROCr/HSA 分配真实内存或建立 SVM range - 把返回的裸指针登记到 runtime 的内存对象表HIP 内存相关 API 很多本文先把它们按语义分组类型HIP API典型语义Device memoryhipMalloc/hipFree分配和释放 GPU device memoryManaged memoryhipMallocManaged/hipMemAdvise/hipMemPrefetchAsync分配和控制 HMM/SVM 管理内存Host memoryhipHostMalloc/hipHostRegisterpinned host memory / registered memoryCopy pathhipMemcpy/hipMemcpyAsync不一定分配但依赖内存对象查询Memory queryhipPointerGetAttributes/hipMemGetInfo查询 runtime 元数据或后端状态Memory poolhipMemPool*/ async mallocHIP runtime 侧的池化分配路径本文主线先关注这些接口hipMalloc hipFree hipMallocManaged hipMemAdvise hipMemPrefetchAsync hipHostMalloc / hipHostRegister其中hipMemAdvise的完整属性翻译链放在了下一篇HIP hipMemAdvise 到 libksamkt中的 hsaKmtSVMSetAttr 调用链分析本文只把它放到内存分配体系里说明它和 SVM range、libhsakmt 的关系。2. 前置背景Topology 发现之后有什么内存分配发生在 HIP runtime 初始化之后。前一篇 topology 文档已经说明初始化后 HIP 层具备这些对象hip::g_devices[N] - hip::Device - amd::Context - amd::Device - roc::Device - hsa_agent_t - HSA memory pools - CPU agent / NUMA 信息 - P2P/link 信息这一步对内存分配很重要因为hipMalloc后续需要回答几个问题当前线程的 current device 是哪个hip::Device这个hip::Device对应哪个amd::Context这个 context 里的amd::Device在 ROCm 后端下是哪一个roc::Device这个roc::Device有哪些 GPU memory pool、fine-grain pool、CPU agent、NUMA 信息如果分配的是 managed/host memory是否需要 HMM/SVM 属性初始化所以 topology 发现结束后不只是“知道有几个 GPU”而是为后续内存分配准备好了当前 HIP device - ROCclr context - ROCclr device - ROCm backend device - HSA agent memory pool3. HIP Runtime API 入口模式HIP 内存 API 入口一般分成两层公共 HIP API - HIP_INIT_API / 参数检查 / stream capture 检查 - ihip* 内部实现以hipMalloc为例hipError_thipMalloc(void**ptr,size_t sizeBytes){HIP_INIT_API(hipMalloc,ptr,sizeBytes);CHECK_STREAM_CAPTURE_SUPPORTED();HIP_RETURN_DURATION(ihipMalloc(ptr,sizeBytes,0),ReturnPtrValue(ptr));}HIP_INIT_API的意义是如果 runtime 还没有初始化就先触发hip::init()完成 device/topology/context 初始化。真正分配逻辑在ihipMalloc()。hipFree也类似hipError_thipFree(void*ptr){HIP_INIT_API(hipFree,ptr);CHECK_STREAM_CAPTURE_SUPPORTED();HIP_RETURN(ihipFree(ptr));}所以调试时公共 API 适合确认用户参数和初始化是否发生ihip*函数才是看业务逻辑的主要入口。4. Device Memory 主路径hipMallochipMalloc的核心实现是ihipMalloc()。简化后的逻辑是hipError_tihipMalloc(void**ptr,size_t sizeBytes,unsignedintflags){booluseHostDevice(flagsCL_MEM_SVM_FINE_GRAIN_BUFFER)!0;amd::Context*curDevContexthip::getCurrentDevice()-asContext();amd::Context*amdContextuseHostDevice?hip::host_context:curDevContext;constautodev_infoamdContext-devices()[0]-info();hip::getCurrentDevice()-SetActiveStatus();*ptramd::SvmBuffer::malloc(*amdContext,flags,sizeBytes,dev_info.memBaseAddrAlign_,useHostDevice?curDevContext-svmDevices()[0]:nullptr);amd::Memory*memObjgetMemoryObject(hip::getCurrentDevice(),*ptr,offset);memObj-getUserData().deviceIdhip::getCurrentDevice()-deviceId();returnhipSuccess;}这里有几个关键点。第一hipMalloc默认使用当前 device 的 per-device contexthip::getCurrentDevice() - asContext() - 当前 device 对应的 amd::Context第二真正创建内存对象的入口是 ROCclr 的amd::SvmBuffer::malloc(...)SvmBuffer::malloc()再转到 contextvoid*SvmBuffer::malloc(Contextcontext,cl_svm_mem_flags flags,size_t size,size_t alignment,constamd::Device*curDev,void*hostptr){void*retcontext.svmAlloc(size,alignment,flags,curDev,hostptr);Add(ret_u,ret_usize);returnret;}第三context.svmAlloc()会根据 context 里的 device 选择后端实现。ROCm 后端是amd::Context::svmAlloc - amd::Device::svmAlloc 虚接口 - roc::Device::svmAllocroc::Device::svmAlloc()会创建一个隐藏的 ROCclr buffermemnew(context)amd::Buffer(context,flags,size,svmPtrUsed);mem-create(nullptr);amd::MemObjMap::AddMemObj(mem-getSvmPtr(),mem);returnmem-getSvmPtr();到这里HIP 用户拿到的是void*但 runtime 内部已经有了用户指针 ptr - amd::MemObjMap - amd::Memory / amd::Buffer - roc::Memory - HSA allocation4.1roc::Memory如何走到 HSA memory poolamd::Buffer::create()会为具体 device 创建后端 memory。ROCm 后端对应roc::Memory::Buffer::create()。对于普通 device local memory最终会走到deviceMemory_dev().deviceLocalAlloc(size(),flags);roc::Device::deviceLocalAlloc()会选择一个 HSA memory poolconsthsa_amd_memory_pool_tpoolflags.pseudo_fine_grain_gpu_ext_fine_grained_segment_.handle?gpu_ext_fine_grained_segment_:flags.atomics_gpu_fine_grained_segment_.handle?gpu_fine_grained_segment_:gpuvm_segment_;然后调用 ROCr/HSA 扩展 APIhsa_status_t statHsa::memory_pool_allocate(pool,size,hsa_mem_flags,ptr);Hsa::memory_pool_allocate()是 ROCclr 对动态加载 ROCr 符号的包装实际符号是hsa_amd_memory_pool_allocate因此hipMalloc的 device memory 主路径可以写成hipMalloc - ihipMalloc - amd::SvmBuffer::malloc - amd::Context::svmAlloc - roc::Device::svmAlloc - amd::Buffer::create - roc::Memory::Buffer::create - roc::Device::deviceLocalAlloc - Hsa::memory_pool_allocate - hsa_amd_memory_pool_allocate - ROCr Runtime::AllocateMemory - hsaKmtAllocMemory - hsaKmtMapMemoryToGPUNodes - KFD5.amd::Memory和内存对象登记HIP API 对外暴露的是裸指针但 runtime 必须知道这个指针对应的内存类型、大小、context、device、flags 等信息。这个信息由amd::Memory保存。分配完成后ROCm 后端会把内存对象登记到全局映射表amd::MemObjMap::AddMemObj(mem-getSvmPtr(),mem);后续很多 API 都会通过用户指针反查amd::Memory*memObjgetMemoryObject(hip::getCurrentDevice(),ptr,offset);这个设计解释了几个现象。第一hipFree(ptr)不能直接释放裸指针它必须先找到amd::Memoryptr-getMemoryObject-amd::Memory-判断是否来自 memory pool/SVM/external memory-选择正确释放路径第二hipMemcpy也需要查内存对象因为它要判断 src/dst 是 host、device、managed、registered host memory还是普通 CPU 指针。第三hipPointerGetAttributes、hipMemRangeGetAttribute、hipMemAdvise这类 API 本质上都是围绕这个 runtime metadata 工作。所以amd::Memory是 HIP 指针式 API 和 ROCclr 对象式内存模型之间的桥HIP 用户模型void* ptr ROCclr 内部模型amd::Memory / roc::Memory 底层驱动模型HSA allocation / SVM range / KFD BO6. Managed Memory / SVM 路径hipMallocManaged不走普通 per-device context而是使用hip::host_contextamd::Contextctx*hip::host_context;constamd::Devicedev*ctx.devices()[0];*ptramd::SvmBuffer::malloc(ctx,CL_MEM_SVM_FINE_GRAIN_BUFFER|CL_MEM_ALLOC_HOST_PTR,size,dev.info().memBaseAddrAlign_);这里的关键点是managed memory 需要表达“系统内存 / SVM / 多 GPU 可访问 / 可迁移”的语义因此它使用包含所有可见 GPU 的host_context而不是某个hip::Device的单设备 context。简化路径是hipMallocManaged - ihipMallocManaged - hip::host_context - amd::SvmBuffer::malloc - amd::Context::svmAlloc - roc::Device::svmAlloc - amd::Buffer / roc::Memory - HMM/SVM 初始化 - Hsa::svm_attributes_set - hsa_amd_svm_attributes_set - ROCr Runtime::SetSvmAttrib - hsaKmtSVMSetAttr - AMDKFD_IOC_SVM在 HMM 支持路径下ROCm 后端会对系统内存做 SVM 初始化if(dev().info().hmmSupported_){deviceMemory_dev().reserveMemory(size(),amd::Os::pageSize());dev().SvmAllocInit(deviceMemory_,size());}SvmAllocInit()内部会调用roc::Device::SetSvmAttributesInt(..., first_alloc true) - Hsa::svm_attributes_set - hsa_amd_svm_attributes_set这一步不是普通的 VRAM allocation而是给一段 SVM/HMM 地址范围建立初始访问属性例如让 GPU agent 可以访问这段系统内存。6.1hipMemAdvisehipMemAdvise不负责分配新内存但它会修改已有 SVM range 的属性。它的核心路径是hipMemAdvise / hipMemAdvise_v2 - ihipMemAdvise - getMemoryObject - amd::Device::SetSvmAttributes - roc::Device::SetSvmAttributesInt - Hsa::svm_attributes_set - hsa_amd_svm_attributes_set - ROCr Runtime::SetSvmAttrib - hsaKmtSVMSetAttr - AMDKFD_IOC_SVM属性会经历几次翻译hipMemoryAdvise - amd::MemoryAdvice - HSA_AMD_SVM_ATTRIB_* - HSA_SVM_ATTR_* / HSA_SVM_FLAG_* - KFD_IOCTL_SVM_ATTR_*这条路径和 managed memory 关系很近因为它操作的是同一类 SVM/HMM 地址范围。6.2hipMemPrefetchAsynchipMemPrefetchAsync也不分配新内存但它会触发 managed/SVM range 向目标位置迁移或建立目标访问倾向。简化路径是hipMemPrefetchAsync - ihipMemPrefetchAsync - getMemoryObject - 解析目标 location / CPU 或 GPU agent - ROCclr command 或 ROCr SVM prefetch/attribute 路径 - hsa_amd_svm_prefetch_async 或 hsaKmtSVMSetAttr(prefetch loc) - KFD SVM range 迁移 / 预取不同 ROCm 版本和 HMM 配置下prefetch 可能表现为显式 prefetch API也可能体现为 SVM 属性更新和迁移请求。调试时要同时关注 ROCr 的hsa_amd_svm_prefetch_async和 libhsakmt 的hsaKmtSVMSetAttr。7. Host Memory 路径hipHostMalloc/hipHostRegisterhipHostMalloc最终也是调用ihipMalloc()但会传入 fine-grain / atomics / NUMA / uncached 等 flagsunsignedintihipFlagsCL_MEM_SVM_FINE_GRAIN_BUFFER;if(flagshipHostMallocUncached){ihipFlags|ROCCLR_MEM_HSA_UNCACHED;}if(flags0||flagshipHostMallocMapped||HIP_HOST_COHERENT){ihipFlags|CL_MEM_SVM_ATOMICS;}if(flagshipHostMallocNumaUser){ihipFlags|CL_MEM_FOLLOW_USER_NUMA_POLICY;}hipError_t statusihipMalloc(ptr,sizeBytes,ihipFlags);因为带有CL_MEM_SVM_FINE_GRAIN_BUFFERihipMalloc()会选择amd::Context*amdContexthip::host_context;所以hipHostMalloc的核心路径是hipHostMalloc - ihipHostMalloc - ihipMalloc(..., CL_MEM_SVM_FINE_GRAIN_BUFFER | ...) - hip::host_context - amd::SvmBuffer::malloc - amd::Context::svmAlloc - roc::Device::svmAlloc - roc::Memory::Buffer::create - hostAlloc / HMM reserve / SVM init - ROCr/libhsakmt/KFDhipHostRegister的语义不同它不是分配一块新 host memory而是把用户已有的 host pointer 注册成 GPU 可访问。它通常会涉及用户已有 host pointer - HIP runtime 创建/登记 memory object - ROCclr/ROCm 后端 lock/register host memory - ROCr HSA memory_register / memory_lock / SVM attributes - libhsakmt pin/map userptr - KFD 建立 GPU 可访问映射因此hipHostMalloc更像“runtime 分配 pinned/SVM host memory”hipHostRegister更像“把已有 CPU 地址范围纳入 GPU 可访问管理”。8. Free 路径hipFreehipFree的关键不是直接释放ptr而是先通过ptr找到 runtime 内部的amd::MemoryhipError_tihipFree(void*ptr){if(ptrnullptr){returnhipSuccess;}amd::Memory*memory_objectgetMemoryObject(hip::getCurrentDevice(),ptr,offset);if(memory_object!nullptr){autodevice_idmemory_object-getUserData().deviceId;if(!g_devices[device_id]-FreeMemory(memory_object,nullptr)){g_devices[device_id]-SyncAllStreams();if(memory_object-getSvmPtr()nullptr){amd::MemObjMap::RemoveMemObj(ptr);memory_object-release();}else{amd::SvmBuffer::free(memory_object-getContext(),ptr);}}returnhipSuccess;}returnhipErrorInvalidValue;}释放路径大致是hipFree - ihipFree - getMemoryObject - 判断是否属于 HIP memory pool - 必要时同步相关 stream - amd::SvmBuffer::free / amd::Memory::release - amd::Context::svmFree - roc::Device / roc::Memory free - hsa_amd_memory_pool_free / hsa_memory_free - hsaKmtFreeMemory - KFD这里要注意释放时使用的是 allocation 时记录在memObj-getUserData().deviceId里的 device而不是简单使用当前线程 current device。这是为了避免用户切换 device 后释放指针时找错设备上下文。9. Copy / Prefetch 为什么也依赖内存对象hipMemcpy本身不是分配接口但它强依赖内存对象表。原因是hipMemcpyDefault、managed memory、host registered memory、device pointer 都要通过 runtime metadata 判断。典型逻辑是hipMemcpyAsync - 解析 stream - getMemoryObjectPairs(src, dst) - 判断 src/dst 分别是什么内存 - 创建 ROCclr copy command - 提交到 hip::Stream / amd::HostQueue - ROCm 后端生成 SDMA 或 kernel/blit 命令也就是说分配阶段建立的ptr - amd::Memory映射会直接影响后续 copy 路径hipMalloc / hipHostMalloc / hipMallocManaged - MemObjMap 登记 - hipMemcpy / hipFree / hipMemAdvise / hipPointerGetAttributes 反查hipMemPrefetchAsync更接近 managed memory 的控制路径它先通过getMemoryObject找到 SVM allocation再根据目标 location 决定是迁移到 GPU、迁移到 CPU还是设置 NUMA/agent 相关属性。10. 到 libhsakmt 的边界在哪里这一节重点看 ROCr Runtime 进入 libhsakmt 的边界上层已经完成 HIP/ROCclr 对象转换之后底层会根据内存类型进入 allocation/map/free 或 SVM attribute 路径。不同 HIP 语义到 libhsakmt 的落点大致如下HIP 语义ROCclr/ROCr 落点libhsakmt/KFD 可能涉及hipMallocdevice memoryhsa_amd_memory_pool_allocatehsaKmtAllocMemory、hsaKmtMapMemoryToGPUNodeshipFreehsa_amd_memory_pool_free/hsa_memory_freehsaKmtFreeMemoryhipMallocManagedSVM/HMM allocation hsa_amd_svm_attributes_sethsaKmtSVMSetAttr、AMDKFD_IOC_SVMhipHostMallocfine-grain/SVM host allocationhsaKmtAllocMemory、map/pin、SVM attributeshipHostRegisterhost memory register/lockuserptr pin/map、GPU mappinghipMemAdvisehsa_amd_svm_attributes_sethsaKmtSVMSetAttrhipMemPrefetchAsyncSVM prefetch / prefetch location attributehsaKmtSVMSetAttr或 prefetch ioctl 相关路径以 device memory 为例ROCr 的hsa_amd_memory_pool_allocate()最终调用core::Runtime::AllocateMemory-hsaKmtAllocMemory-hsaKmtMapMemoryToGPUNodeslibhsakmt 的 memory path 会进入libhsakmt/src/memory.c-hsaKmtAllocMemoryCtx-hsaKmtAllocMemoryAlignCtx-ioctl(KFD allocation/map memory path)以 SVM 属性为例ROCr 的hsa_amd_svm_attributes_set()最终调用core::Runtime::SetSvmAttrib-hsaKmtSVMSetAttr-libhsakmt/src/svm.c-hsakmt_ioctl(AMDKFD_IOC_SVM)所以如果目标是观察“HIP 内存 API 什么时候到 libhsakmt”断点不应该只打在 HIP 层还要打在 ROCr 和 libhsakmt 边界。11. 调试建议可以按层级设置断点。HIP 层hipMalloc hipFree hipHostMalloc hipMallocManaged hipMemAdvise hipMemPrefetchAsync ihipMalloc ihipFree ihipHostMalloc ihipMallocManaged ihipMemAdvise ihipMemPrefetchAsync getMemoryObject getMemoryObjectPairsROCclr 层amd::SvmBuffer::malloc amd::SvmBuffer::free amd::Context::svmAlloc amd::Context::svmFree amd::Buffer::create amd::MemObjMap::AddMemObj amd::MemObjMap::FindMemObj roc::Device::svmAlloc roc::Device::deviceLocalAlloc roc::Device::hostAlloc roc::Device::SetSvmAttributes roc::Device::SetSvmAttributesInt roc::Device::SvmAllocInit roc::Memory::Buffer::createROCr 层hsa_amd_memory_pool_allocate hsa_amd_memory_pool_free hsa_amd_svm_attributes_set hsa_amd_svm_attributes_get hsa_amd_svm_prefetch_async core::Runtime::AllocateMemory core::Runtime::SetSvmAttriblibhsakmt 层hsaKmtAllocMemory hsaKmtAllocMemoryAlign hsaKmtMapMemoryToGPU hsaKmtMapMemoryToGPUNodes hsaKmtFreeMemory hsaKmtSVMSetAttr hsaKmtSVMGetAttr hsakmt_ioctl如果只想看hipMalloc主路径推荐第一轮断点是hipMalloc - ihipMalloc - amd::SvmBuffer::malloc - roc::Device::svmAlloc - roc::Memory::Buffer::create - roc::Device::deviceLocalAlloc - hsa_amd_memory_pool_allocate - core::Runtime::AllocateMemory - hsaKmtAllocMemory - hsaKmtMapMemoryToGPUNodes如果只想看 managed/SVM 路径推荐第一轮断点是hipMallocManaged - ihipMallocManaged - amd::SvmBuffer::malloc - roc::Device::svmAlloc - roc::Device::SvmAllocInit - roc::Device::SetSvmAttributesInt - hsa_amd_svm_attributes_set - core::Runtime::SetSvmAttrib - hsaKmtSVMSetAttr - hsakmt_ioctl(AMDKFD_IOC_SVM)12. 总结HIP 内存分配接口到 libhsakmt 的链路可以压缩成一张图HIP API - HIP runtime object model - ROCclr context/device/memory model - ROCm backend roc::Device / roc::Memory - ROCr HSA memory/SVM API - libhsakmt thunk - KFD ioctl其中最重要的是两个对象桥接hip::Device - amd::Context - amd::Device / roc::Device和void* ptr - amd::MemObjMap - amd::Memory / roc::Memory前者决定“这次分配属于哪个 device/context”后者决定“后续拿着这个裸指针还能不能找回 runtime 内部的内存对象”。所以hipMalloc到 libhsakmt 不是一条简单的函数转发链而是一条分层资源建模链HIP 层提供 CUDA-like 指针 API。ROCclr 层创建 context-scoped memory object。ROCm 后端选择 HSA memory pool 或 SVM/HMM 路径。ROCr Runtime 翻译成底层 allocation / map / SVM attribute 操作。libhsakmt 把这些操作提交给 KFD。理解这条链路后再看hipFree、hipMemcpy、hipMemAdvise、hipMemPrefetchAsync就会顺很多它们都是围绕同一个amd::Memory元数据和同一组 ROCr/libhsakmt 底层能力展开。