本篇讲清楚用户态如何触发 / 控制 UVM 迁移尤其是回迁到 CPU 时如何指定 NUMA 节点。分三个层面CUDA 高层 API、UVM 用户库 API、底层 ioctl 协议。1. 三层接口关系应用 / CUDA Runtime │ cudaMemPrefetchAsync / cudaMemAdvise / cudaMallocManaged ▼ libnvidia-uvm 用户库 (UvmMigrate / UvmSetPreferredLocation ...) │ ioctl(/dev/nvidia-uvm, UVM_MIGRATE, ...) ▼ UVM 内核驱动 (uvm_api_migrate, uvm_api_set_preferred_location ...)开源仓库kernel-open/nvidia-uvm/uvm_ioctl.h定义了内核 ABICUDA/UVM 用户库是这些 ioctl 的封装。下面以 ioctl 结构体为准这是可复现的 ground truth。2. 主动迁移UVM_MIGRATE定义见 uvm_ioctl.h#defineUVM_MIGRATEUVM_IOCTL_BASE(51)typedefstruct{NvU64 base;// IN 起始虚拟地址NvU64 length;// IN 字节长度NvProcessorUuid destinationUuid;// IN 目标处理器 UUIDGPU 的 UUIDCPU 用特殊 UUIDNvU32 flags;// IN UVM_MIGRATE_FLAG_*NvU64 semaphoreAddress;// IN 异步完成信号可选NvU32 semaphorePayload;// INNvS32 cpuNumaNode;// IN ★ 目标 CPU NUMA 节点回迁到 CPU 时使用NvU64 userSpaceStart;// OUT 需用户态补做的区间起点NvU64 userSpaceLength;// OUT 需用户态补做的区间长度NV_STATUS rmStatus;// OUT}UVM_MIGRATE_PARAMS;2.1cpuNumaNode语义关键当目标是 CPU 时cpuNumaNode指定希望页面落到哪个 NUMA 节点。合法性校验见 uvm_ioctl.h 注释 与uvm_api_migrate中的校验uvm_migrate.ccpuNumaNode被视为非法若小于-1大于等于系统最大节点数对应一张已注册的 GPU 的内存节点不在node_possible_map内该节点没有 online 的内存!nv_numa_node_has_memory。特殊值cpuNumaNode -1 (NUMA_NO_NODE)对managed 内存表示「不限定节点交给内核/策略决定」对pageable 内存则是非法pageable 路径必须给出确定节点见下文重试协议。2.2 flagsflag含义UVM_MIGRATE_FLAG_ASYNC异步配合semaphoreAddress/Payload完成通知UVM_MIGRATE_FLAG_SKIP_CPU_MAP目标为 CPU 时跳过建立 CPU 映射仅测试构建可用使迁移可完全异步UVM_MIGRATE_FLAG_NO_GPU_VA_SPACE允许目标 GPU 未注册 GPU VA space仅测试2.3 Pageable 内存的「用户态-内核态协作」重试协议对系统分配pageable内存一次 ioctl 未必能完成用户库需要按内核返回码循环处理见 uvm_ioctl.h 注释NV_WARN_NOTHING_TO_DO内核遇到 file-backed vma 或无 GPU 可驱动拷贝。用户态改用move_pages(2)迁移userSpaceStart/userSpaceLength指示的区间然后从该 vma 之后继续。NV_ERR_MORE_PROCESSING_REQUIRED在目标 CPU 节点分配失败。用户态换一个 CPU NUMA 节点遵循线程的 NUMA 策略重试若无更多节点可试则改用UVM_POPULATE_PAGEABLE在任意节点补页。NV_OK成功仅保证页面被 populate不保证一定落在请求节点。这段协议正是「回迁到 CPU 时考虑 NUMA」在用户态的体现内核会尽力落在cpuNumaNode落不下就把决策权交回用户态换节点重试。3. 首选位置策略UVM_SET_PREFERRED_LOCATION定义见 uvm_ioctl.h#defineUVM_SET_PREFERRED_LOCATIONUVM_IOCTL_BASE(42)typedefstruct{NvU64 requestedBase;// INNvU64 length;// INNvProcessorUuid preferredLocation;// IN 首选处理器CPU 或某 GPUNvS32 preferredCpuNumaNode;// IN ★ 首选位置为 CPU 时的首选 NUMA 节点NV_STATUS rmStatus;// OUT}UVM_SET_PREFERRED_LOCATION_PARAMS;作用为一段 VA range 设置「首选驻留位置」。之后任何原因触发的「回迁到 CPU」——无论是显式UVM_MIGRATE(cpuNumaNode-1)、缺页、还是驱逐——只要没有更高优先级的显式节点就会优先落到preferredCpuNumaNode。内核侧写入policy-preferred_location与policy-preferred_nidmanageduvm_va_range_set_preferred_location()uvm_va_range.cHMM/系统内存uvm_va_policy_set_preferred_location()uvm_va_policy.c。对应 CUDAcudaMemAdvise(ptr, size, cudaMemAdviseSetPreferredLocation, device)CPU 节点粒度的首选位置对应较新的cudaMemLocationnode 类型语义。UVM_UNSET_PREFERRED_LOCATIONIOCTL 43清除策略preferred_nid回到NUMA_NO_NODE。4. 查询当前策略UVM_TOOLS_GET_PROCESSOR_UUID_TABLE/ 通过uvm_test路径可读回策略。内核在导出参数时把policy-preferred_nid写入params-preferred_cpu_nidmanageduvm_va_range.c默认NUMA_NO_NODE有策略时填policy-preferred_nidHMMuvm_hmm.c。5. 典型用户态用法示例5.1 CUDA 高层// 分配 managed 内存void*p;cudaMallocManaged(p,N);// 建议首选落在 CPU 的 NUMA node 2较新 CUDA 用 cudaMemLocation 表达 nodecudaMemLocation loc{.typecudaMemLocationTypeHostNuma,.id2};cudaMemAdvise_v2(p,N,cudaMemAdviseSetPreferredLocation,loc);// 在 GPU 上算完后把数据预取回 CPU node 2cudaMemLocation host{.typecudaMemLocationTypeHostNuma,.id2};cudaMemPrefetchAsync_v2(p,N,host,/*flags*/0,stream);cudaStreamSynchronize(stream);// 此时 p 的页面会尽量驻留在 node 2 的 DRAM 上5.2 直接走 ioctl做驱动/系统软件时intfdopen(/dev/nvidia-uvm,O_RDWR);// 回迁 [base, baselen) 到 CPU 的 NUMA node 1UVM_MIGRATE_PARAMS params{0};params.basebase;params.lengthlen;params.destinationUuidUVM_CPU_UUID;// CPU 的约定 UUIDparams.cpuNumaNode1;// ★ 目标 NUMA 节点params.flags0;// 同步for(;;){ioctl(fd,UVM_MIGRATE,params);if(params.rmStatusNV_OK)break;elseif(params.rmStatusNV_WARN_NOTHING_TO_DO){// 用 move_pages() 处理 [userSpaceStart, userSpaceLength)再从其后继续move_pages_range(params.userSpaceStart,params.userSpaceLength,/*node*/1);params.baseparams.userSpaceStartparams.userSpaceLength;params.lengthoriginal_end-params.base;}elseif(params.rmStatusNV_ERR_MORE_PROCESSING_REQUIRED){params.cpuNumaNodepick_another_numa_node();// 换节点重试params.baseparams.userSpaceStart;}else{// 真正的错误break;}}注意UVM_CPU_UUID是 UVM 约定用于表示「CPU 处理器」的特殊 UUID实际值以用户库/头文件为准。6. 用户态需要记住的 NUMA 要点两个入口都能带节点一次性迁移用UVM_MIGRATE.cpuNumaNode长期倾向用UVM_SET_PREFERRED_LOCATION.preferredCpuNumaNode。优先级显式cpuNumaNodepreferred_nid 内核默认。managed 允许-1不限定pageable 不允许-1。不保证严格落点内核尽力__GFP_THISNODE失败会回退到任意节点或把控制权交回用户态pageable 的重试协议。Grace-Hopper / EGM 等迁到「集成 GPU」实际等价于迁到该 GPU 最近的 CPU 节点见 04-special-hardware-ats。内核侧如何消费这些参数见 02-kernel-migration-flow 与 03-numa-node-selection。
