GPU 访问页表不存在时的完整处理流程这名字听起来像是纯驱动开发人员才会碰的问题但实际上只要你在 GPU 上写过 CUDA、OpenCL 或者用过统一内存Unified Memory你就迟早会遇到它。我最早对这个流程产生兴趣是因为一次非常诡异的崩溃程序跑一两个小时才崩一次错误信息只有一条CUDA error: an illegal memory access was encountered没有越界、没有空指针所有代码看起来都是对的。后来一步步排查才发现根子出在 GPU 访问一个“不该存在”的页表项上。这篇文章想把它彻底讲清楚GPU 什么时候会访问到不存在的页表项、硬件层面是怎么发现这件事的、驱动收到通知之后会做什么、以及不同厂商的 GPU 在这个环节上有哪些差异。内容会把硬件的处理机制和驱动侧的逻辑分开来拆最后给出一套实际项目里能用的排查和规避方法。1. 为什么 GPU 会踩到一个“不存在”的页表项1.1 先从 GPU 的虚拟内存模型说起GPU 和 CPU 一样对显存和系统内存的访问都走“虚拟地址 - 页表 - 物理地址”的翻译路径。GPU 的地址空间会被拆成固定大小的页常见的页大小是 4KB、64KB有些新架构还支持 2MB 甚至 1GB 的大页。每一级页表负责把虚拟地址映射到下一级页表或者最终映射到物理页框。页表项的状态一般可以分成三类有效Valid并且映射到物理内存、无效Invalid但属于某个合法分配范围、完全不存在Not Present。前两种在驱动里都有对应的管理数据结构第三种才是最危险的——GPU 去读页表的时候发现这一项连“被分配过”的标记都没有。注意GPU 缺页和 CPU 缺页在语义上有本质区别。CPU 缺页时操作系统可以通过页面换入page-in让程序无感继续执行GPU 缺页时能不能恢复、如何恢复完全取决于驱动是否实现了对应的机制。没有实现机制的驱动任何缺页都会被当作 fatal error 处理。1.2 哪些典型场景会触发 GPU 访问不存在的页表项第一种是赤裸裸的越界访问。比如你分配了一个长度为N的数组但 kernel 里的索引因为算法边界没算对跑到了N 偏移的位置。这个地址大概率落在分配范围之外对 GPU MMU 来说就是“页表项不存在”。这类问题在调试时最容易定位但在跑大规模并行任务时线程号、块号、步长任何一个算错都会产生这种访问。第二种和设备内存的惰性分配lazy allocation有关。GPU 驱动为了省内存在cudaMalloc或者clCreateBuffer之后可能并不会立刻把物理页真正挂到页表里。只有在 kernel 首次访问对应地址时驱动才借缺页的时机去完善映射。此时如果驱动在缺页处理中发现用户程序访问的区域已经越过了分配的边界这个缺页就会升级为错误。第三种和统一内存Unified Memory/ 共享虚拟内存SVM相关。在这种模型下CPU 和 GPU 共享一个虚拟地址空间物理页面可以在系统内存和显存之间搬移。程序访问一个此前从未触达过的地址时页表里确实没有这一项驱动需要走一套完整流程把物理页面找回来。第四种是多 GPU 场景下的 Peer-to-Peer 访问。当你让 GPU A 直接访问 GPU B 的显存时如果 A 的页表里没有 B 那边物理内存的映射A 会发出一个跨设备访问B 的 MMU 会返回一个错误给 A。这里任何一方没有处理好最终都会表现为“访问了不存在的页表项”。1.3 缺页不一定是 bug也可能是特性我这里要强调一下GPU 缺页并不总是坏事。在支持完整缺页处理流程的平台上“缺页时按需填充页表”是一种节省内存、提升启动速度的正面机制。CUDA 的 Managed Memory 之所以能支持比物理显存更大的工作集靠的就是这套处理流程。问题只在驱动有没有把流程走完而不是缺页本身有多可怕。2. 硬件层的“哨兵”GPU MMU 怎么发现页表项缺失2.1 从 TLB 命中失败到页表遍历现代 GPU 都内置了 TLBTranslation Lookaside Buffer用来缓存最近的虚拟地址到物理地址的翻译结果。GPU 的核心要访问一个地址时第一步是查 TLB命中就直接生成物理地址没命中就要走到多级页表里去逐级查找。页表遍历的过程在硬件上是由 GPU MMU 完成的。它会先找到顶级页表的基地址然后逐级往下走每一级根据虚拟地址里对应的位段去索引页表项。如果某一级页表项被标记为 “Not Present”遍历就会中止不会继续往下走。这里有一个需要留意的细节页表遍历自身的缓存Page Walk Cache可能让硬件在缺页之后继续读到旧数据。驱动在处理完缺页、更新完页表之后不能只依赖 TLB 刷新的命令还要确保 PCIe ATS 相关的缓存和页表遍历缓存一并失效。很多“CPU 能看见新页表、GPU 仍然报错”的诡异问题根子就在这里。2.2 GPU 的缺页通知机制中断、陷阱还是错误日志不同 GPU 厂商在“发现页表项缺失后如何上报”上走了完全不同的路。NVIDIA 的方案是把缺页信息写进 GPU 的错误寄存器并给驱动发出中断。驱动通过读取寄存器拿到出错的虚拟地址、访问类型读/写/原子操作以及请求的进程上下文 ID。CUDA 运行时碰到这种情况如果缺页发生在驱动没有准备恢复的路径上就会向上返回illegal memory access。AMD 在 GFX9 及之后的架构里设计了更细粒度的 VM fault 处理机制。硬件会记录 fault 地址、fault 类型并通过中断通知驱动。AMD 的 GPUVM 模块把 VM fault 分成了可恢复和不可恢复两种可恢复的 fault 会走迁移流程不可恢复的会直接给应用返回 GPU hang 或者 device lost。Intel 的新架构Arc 系列走的是 Shared Virtual Memory 路线缺页处理整个嵌进了 Linux 内核的 HMMHeterogeneous Memory Management框架。GPU 缺页会被转换成类似 CPU 缺页的语义走 mmu_notifier 一整套机制。2.3 GPU 执行状态在缺页瞬间会发生什么缺页发生在 GPU 某个执行单元正在处理一条内存访问指令的时候。这个执行单元不能像什么也没发生一样继续执行下一条指令因为它没有拿到合法的物理地址。所以硬件必须让触发缺页的 wavefront、warp 或 workgroup 进入等待状态。在部分 NVIDIA 架构上缺页会导致整个 GPU context 被阻塞。也就是说不只出错的线程停下来整个 context 里的所有线程都一起停等缺页流程处理完再统一恢复。这种“全堵”的机制简化了恢复逻辑但也让缺页的代价被放大一次缺页可能让整卡吞吐瞬间掉到接近零。AMD 的实现更接近“局部阻塞”。硬件尝试只停掉发起 fault 的 wavefront其他不相干的 wavefront 可以继续执行。但在驱动需要刷新页表、做全局 TLB shootdown 的情况下局部恢复很难做到严格隔离最后还是会出现一定范围的停顿。3. 驱动接管缺页后的完整处理链路3.1 第一站确认地址是否属于合法分配GPU 硬件发出缺页通知后驱动要做的第一件事不是立刻分配物理页面而是确认这个虚拟地址在“谁”的分配范围里。这是整个流程里最关键的分叉点地址合法走恢复路径地址非法直接报错并杀死或标记相关任务。驱动维护了一张虚拟地址区间表类似 CPU 的 VMAVirtual Memory Area。在 CUDA 语境里每个cudaMalloc返回的地址区间都会被记录在这张表里。驱动拿到缺页地址后会按区间树的查找逻辑去判断这个地址落在哪个区间内这个区间的类型是什么普通显存、Managed Memory、Peer 映射、主机内存映射我见过一个很隐蔽的问题地址落在某个合法的分配区间内但页表的生成过程因为并发问题把这个区间对应的页表项给漏掉了。程序“访问了不该访问的页表项”但驱动按地址合法性检查却能查到区间这时会让驱动误以为是可以恢复的缺页。走完恢复流程之后会发现要么映射恢复成功要么因为区间本身没有申请物理页面而走二次错误路径。3.2 恢复路径之一按需物理页面分配对于合法的普通显存分配缺页恢复的下一步是分配物理页面。驱动从显存分配器里拿一个空闲物理页框然后根据缺页地址计算出对应的页表项索引把物理页框写入页表项并标记为有效。如果缺页地址属于大页区间驱动还需要先把中间层级的页表项创建出来。比如 GPU 用的是两级页表而目标区间被标记为 2MB 大页那么驱动要先确保顶级页表里有指向二级页表的有效项再在二级页表中写入大页映射。3.3 恢复路径之二Managed 内存的迁移决策如果缺页地址属于 Managed Memory / UVM 区间事情就不只是分配物理页面那么简单了。GPU 显存里可能没有这块数据数据当前可能在系统内存里甚至可能在另一块 GPU 的显存里。驱动的处理流程大致这样先锁定这个虚拟地址对应的物理页面pin 住防止迁移过程中 CPU 端换页根据当前数据所在位置决定迁移目标如果是 CPU 端迁往 GPU就发起一个 DMA copy数据到达显存后把对应的显卡页表项指向这份数据副本最后向 CPU 的内存管理器发送通知让它在需要时再进行反向迁移。这套流程在 CUDA 里对应cudaMemPrefetchAsync的手动优化以及cudaMemAdvise提供的访问策略提示。如果程序不做任何提示驱动只能按“访问过的就往显存放、显存压力大了就往系统内存赶”这种被动策略来兜底缺页的代价会明显偏高。3.4 杀招路径非法访问如何结束任务当驱动确认缺页地址不属于任何合法分配区间或者这个区间的访问权限不允许当前执行的操作时恢复流程就中止了进入错误处理。驱动会给对应的执行流stream或 context 打上错误标记丢弃或延迟当前 kernel 的结果。对用户来说最直观的表现是cudaDeviceSynchronize()返回错误。但实际上错误的发生时间比这个早得多kernel 执行过程中缺页就已经发生了后续整个 context 的执行都被冻结直到同步点才把错误返回给 CPU 侧。这也是为什么非法访问出错后后续 kernel 也可能“莫名其妙”地跟着失败——因为它们都在同一个被污染的 context 里。3.5 页表更新之后TLB 失效与一致性同步页表项被修改后GPU 的 TLB 和页表遍历缓存里还留着旧的“不存在”翻译结果必须显式刷新。GPU 的 TLB 刷新范围和 CPU 类似可以按虚拟地址、按地址空间或者全局来刷。刷新动作本身有巨大开销。全局 TLB 刷新意味着 GPU 上所有正在执行的任务都必须在安全点停下来等刷新完成才能恢复。在多进程共享一个 GPU 的场景下一个进程的缺页恢复可能让其他进程的 kernel 跟着抖动。实践中我见过缺页率不高、但整卡吞吐波动剧烈的情况查下来就是 TLB 刷新的全局停顿造成的。重要提示驱动逻辑上必须保证“页表更新”和“TLB 失效”是一个原子序列。先改页表、后失效没问题先失效、后改页表就会出现其他执行单元在 TLB 刷新后仍然能读到旧页表项的时间窗口。这个顺序错误在并发压力下会表现为极难复现的随机错误。4. 驱动侧的关键细节从页表层级到并发安全4.1 GPU 页表的层级设计与大页支持GPU 页表通常是多级的常见设计是类似 CPU 的四级结构但不同厂商往往会砍掉一些层级来简化。层级越少单次缺页恢复需要创建的中间页表项就越少。大页Large Page在 GPU 缺页恢复流程里扮演的角色很特别。如果驱动能使用 2MB 甚至 1GB 的大页做映射一次页表填充就能覆盖很大的地址范围TLB 的命中率也会大幅提升。代价是物理内存分配粒度变大显存碎片问题更难处理。实际项目里对数据量大、访问模式规律的计算任务比如矩阵乘法、卷积我倾向于让驱动开大页对频繁分配、释放的临时缓冲坚持用 4KB 小页反而能减少内存浪费。4.2 并发缺页多个线程同时踩到空页表项大规模并行程序可能在同一时刻让几十个 GPU 线程访问同一个未映射区域。硬件上报给驱动的却可能是多个独立 fault或者一个经过合并的 fault 批次。驱动在处理第一个 fault 时会把页面映射好处理后续 fault 时发现页表已经有效就只需要做一次 TLB 刷新。但如果驱动对每个 fault 都独立走“分配页框 - 更新页表 - 刷 TLB”的流程后果就是重复分配、重复刷新甚至在极端情况下同一地址被映射两次造成物理页泄漏。成熟的驱动会用 per-process 的锁把同一地址的并发 fault 串行化并在页表更新后做一次“重新检查页表是否已有效”的动作。4.3 多 GPU 与 P2P 访问中的特殊处理在多 GPU 系统中驱动为每个 GPU 维护独立的页表结构。GPU A 的页表里要插入一个指向 GPU B 显存地址的映射并不是直接把 B 的物理地址抄进 A 的页表项就行还要考虑 B 那边是否需要反向映射以及 ATSAddress Translation Services在 PCIe 层面如何处理跨设备翻译。这类映射在缺页处理时最麻烦。A 访问 B 的显存发生缺页驱动要先判断 B 上对应的物理页是否存在。如果 B 的显存页由于压力被迁移走了A 的映射就悬空了。恢复的时候实际上要把数据迁回 B 的显存再刷新 A 的页表和 TLB。整个过程涉及的锁和同步点比单 GPU 缺页多得多。4.4 驱动崩溃与设备恢复的兜底机制如果驱动在缺页处理流程里自身出现错误比如页表更新了一半被调度打断或者分配器返回失败GPU 会长时间处于“缺页未处理”的停顿状态。硬件看门狗watchdog超时后驱动会强制重置 GPU并把已经执行的任务标记为失败。这个层面有一个我在项目里踩过的坑GPU 重置会清空显存内容但 CPU 侧的内存复制队列不一定知道这件事。重置发生后后续 DMA 操作还在引用旧的 GPU 地址最终表现为数据全错或者驱动挂起。后来我在所有 DMA 操作前加了设备健康状态检查才把这类问题控制住。5. 平台差异NVIDIA、AMD、Intel 的处理风格5.1 CUDA / UVM用户态更多介入NVIDIA 在 CUDA 的 Managed Memory 里做了非常完整的缺页用户态处理。cudaMemPrefetchAsync本质上是让驱动提前执行“页面迁移 页表更新 TLB 刷新”全套动作把缺页的运行时开销平摊到可预知的时间点上。cudaMemAdvise则允许程序提示驱动某段内存主要被谁访问帮助驱动做迁移决策。如果你不提供任何提示NVIDIA 驱动依靠的是简单的“最近访问”策略。双 GPU 环境下数据在 A、B 两块卡之间来回弹跳每跳一次都可能伴随一次真实的页迁移和两次 TLB 刷新。对这种场景我在实际优化中通常要求程序显式指定数据的“归属 GPU”避免驱动在每次访问时都做迁移判断。5.2 AMD GPUVM更靠近操作系统的模块化AMD 的 Linux 驱动把 GPU 虚拟内存管理集中在内核的amdgpu_vm模块里底层通过 GPUVM 硬件实现页表遍历。AMD 对 HMM 的支持比 NVIDIA 更激进这让 AMD GPU 可以跟 Linux 的mmu_notifier、migrate_vma这些机制深度绑定。好处是CPU 的mmap出来的内存在 GPU 上缺页时驱动的处理路径跟 CPU 缺页非常接近理论上可以实现更细粒度的按需映射。副作用则是调试时需要同时理解 GPU 驱动和内核内存管理两个领域的逻辑报错信息也经常分散在不同的内核日志层里。用 AMD GPU 跑统一内存任务时我会优先查/sys/kernel/debug/amdgpu/里的 VM 相关计数器和内核日志里的 GPUVM fault 条目。5.3 IntelSVM 一体化Intel 的 GPU 驱动走一体化 SVM 路线把 GPU 当作一个普通的内存访问者看待。Intel Arc 系列之后驱动把 GPU 页表直接嵌入 Linux 内核的页表管理框架CPU 和 GPU 的地址空间在操作系统层面统一处理。5.4 平台对比与选型建议平台缺页处理恢复机制主要恢复手段用户态可控程度Linux 内核集成度NVIDIA CUDA按需映射 UVM 迁移锁页、迁移、预取高prefetch/advise中私有 UVM 实现AMD GPUVMHMM 融合的按需映射迁移、页表填充中依赖内核接口高完全内核模块化Intel SVM内核页表融合处理依赖 Linux mm 机制低系统级决定高框架级集成选型上的个人建议如果你的业务跑在 CUDA 生态里优先把用户态的预取和迁移接口用好缺页恢复流程走的是 NVIDIA 私有路径如果你做的是 Linux 内核态加速应用AMD 的 HMM 路线更“标准”跟内核 mm 系统的联动更自然Intel 平台适合不想分别管理 CPU/GPU 内存映射的统一内存场景但目前生态成熟度仍在追赶前两家。6. 实战如何观测、定位和规避 GPU 缺页6.1 缺页是可以量化的很多开发者聊到 GPU 缺页时总觉得它是“随机崩一下”的那种不可观测事件。实际上在多数现代平台上缺页是可以量化的。CUDA 的cudaDeviceGetAttribute拿不到缺页计数但 NVML 和 compute-sanitizer 能提供部分信息。AMD 的内核调试接口里能看到 VM fault 次数。Intel 的perf子系统也在逐步加入 GPU 页表活动的事件。如果只是想要一个粗略的评估可以用一个更土的办法把 kernel 访问的每个内存区域在第一轮访问前手动触发一次全量预取观察预取前后的端到端耗时差。差值越大说明缺页开销在整体中的占比越高。这个方法不精确但快速有效。6.2 用 compute-sanitizer 定位非法访问当缺页是由非法访问引起时NVIDIA 平台的首选工具是 compute-sanitizer。它不是单纯告诉你“有非法访问”而是能定位到具体的 kernel 启动参数、线程坐标、block 坐标。实际上它在 CUDA 12.0 之后已经能报告虚拟地址和页表相关的错误细节。使用方式很简单compute-sanitizer --tool memcheck ./your_program跑出来的报告会指出是“读”还是“写”、访问的地址范围以及哪一行代码发起的访问。这里的地址范围恰好可以和驱动的虚拟地址区间做对照直接判断出缺页地址是否落在合法分配范围内。6.3 读 NVIDIA 的 XID 错误驱动无法恢复的缺页会在系统日志里留下 XID 错误号。比较有代表性的是 XID 31 和 XID 43前者是非法内存访问后者是 GPU 停止响应。结合nvidia-smi -q里的 Ecc 错误计数可以快速判断是否真的发生了页表层面的故障。我处理过一个周期性的“跑 8 小时必崩”案例dmesg里 XID 31 出现前完全没有分配错误或越界提示。后来用cuda-memcheck老工具重新跑了一遍所有测试路径才找到一段仅在特定输入下会越界的代码。这类随机出现的缺页错误最好做二值化搜索缩小输入范围找到首个必现的边界条件。6.4 避免缺页开销的常见策略缺页恢复流程再快也比预先把页表填好慢几个数量级。任何追求性能的 GPU 程序都应该把缺页当“业务逻辑错了”来对待而不是指望缺页处理流程帮你兜底。策略一预热。kernel 启动前用一次空读或者空写把关键数据区域全部触达一遍。CUDA 的cudaMemPrefetchAsync(..., cudaCpuDeviceId)也可以把内存提前迁回 CPU减少跨设备迁移的缺页次数。策略二显存复用。频繁分配/释放显存会导致页表反复变更。用一个内存池把相同大小的缓冲区缓存起来页表项只要建立一次后续 kernel 直接复用。策略三锁页pinned memory。对于跨 CPU 和 GPU 的数据交换区域锁页能让 DMA 路径直接访问物理地址减少缺页处理过程中的二次页面锁定操作。CUDA 里用cudaHostAlloc分配 pinned 内存OpenCL 里用CL_MEM_ALLOC_HOST_PTR配合clEnqueueMapBuffer达到相似效果。策略四控制超额订阅。Managed Memory 的好处是支持超额订阅坏处是频繁换入换出。如果你的工作集经常超过显存容量老老实实做显存预算管理比依赖迁移策略可靠得多。6.5 调试多 GPU 缺页时的专项检查多 GPU 场景下缺页检查要多做几个动作核验 P2P 访问是否真的被驱动允许检查cudaDeviceCanAccessPeer的返回值确认目标地址没有被目标 GPU 迁移走必要时在访问前后各做一次 prefetch留意跨进程共享内存的情况其他进程对页表的修改不会自动同步到本进程的 GPU 地址空间需要用进程间同步机制协调。我在一个双卡推理项目中遇到过 P2P 访问稳定复现缺页的问题排查后发现的根因竟然是目标 GPU 上做了显存压缩压缩后的物理页在 P2P 访问中不满足地址对齐要求。换成非压缩的分配方式后问题立刻消失了。这类问题不一定能在驱动日志里直接看到需要在分配时留意对齐和压缩标志。6.6 最后的小技巧区分“缺页恢复”和“缺页被吞”驱动在高速缺页风暴下可能会选择丢掉部分 page fault 记录只在后续报告中给一个“too many pending faults”之类的汇总。这种情况在 AMD 的部分老驱动上出现过表现为少量随机崩溃但 dmesg 里没有任何完整 fault 信息。遇到这种局面我一般会先看驱动版本再查厂商的 release notes 里有没有修过 fault 风暴相关的补丁最后才去怀疑应用层。之前遇到过一版驱动连续崩溃四次升一个版本后完全稳定就是用了这个排查顺序。7. 写在最后的经验之谈GPU 缺页处理这条链路从硬件 MMU 触发、驱动接管、页表更新到 TLB 失效每一步都有对应的失败模式和优化空间。真正让我觉得“值回票价”的收获并不是背下缺页流程的每一步而是建立了一个习惯把 GPU 内存访问的每一个地址都当作可能触发页表缺失的行为来对待。在这个前提下prefetch、锁页、大页、内存池这些优化手段才会真正被用在该用的地方而不是照着文档盲目加。如果你在自己的项目里遇到了非法访问、随机崩溃或频繁的显存迁移开销建议先不要急着改业务逻辑而是花半小时把日志里的缺页信息、地址区间和分配记录对齐一遍。很多时候问题和业务算法没有关系纯粹是页表状态出了问题。这套排查方式我用了很多年到现在仍然觉得是最稳的起点。