DevIOus 精读：用 GPU 和网卡从 IOMMU 里窥探隔壁设备

Posted on 2026-10-14 Edited on 2026-10-13 Views: Word count in article: 7.4k Reading time ≈ 14 mins.

设备访问内存要先经 IOMMU 做地址翻译，翻译结果缓存在 IOTLB 里，而这块缓存被同一个 IOMMU 下的所有设备共用。于是攻击者不必跟受害者共享任何内存，只要让自己的 GPU 或网卡反复发 DMA 请求并给时延计时，就能从 IOTLB 的状态变化里看出隔壁设备在干什么——跨 VM 成立，隔着 RDMA 也能用。

论文：DevIOus: Device-Driven Side-Channel Attacks on the IOMMU（IEEE S&P 2023）

作者：Taehun Kim、Hyeongjin Park、Seokmin Lee、Junbeom Hur、Youngjoo Shin（Korea University）、Seunghee Shin（SUNY Binghamton）

CPU 缓存侧信道的前提通常写得很清楚：攻击者与受害者要共享内存、共享某一级 CPU 缓存（比如所有核共用的末级缓存）、共享某个地址空间。这篇论文把前提换掉了——不再从 CPU 发起请求，而是让设备自己发。这一换带来两个后果：攻击者不需要与受害者共享任何内存，而且整条信道不经过 CPU 缓存与 TLB。后面那一点，恰好是它能绕开一大堆现成防御的原因。

📚 基本概念速读

名称	定义	记忆点
DMA（Direct Memory Access）	设备不经 CPU 直接读写主存（也就是内存条）	全文所有请求都是它
虚拟地址 / 物理地址	程序与设备用的是虚拟地址，真正落到内存条上的位置叫物理地址	两者之间的换算就是地址翻译
页（page）	内存被切成大小固定的块，翻译和分配都以块为单位	常见 4 KB，大页 2 MB
TLB	CPU 侧缓存地址翻译结果的地方	设备侧的对应物就是 IOTLB
组（set）/ 路（way）	缓存不是一个大池子：地址先决定落在哪一组，每组只有固定几个位置，那些位置叫路	一个组填满后，要按替换策略挤掉旧的
IOMMU	CPU 里管设备 DMA 的硬件，负责把设备的 I/O 地址翻译成物理地址并检查权限	本文的攻击面就在它内部
DRHU	IOMMU 里真正做地址翻译的单元	每次 DMA 请求都要找它
IOVA	设备侧使用的虚拟地址	每个设备有自己的一份
IOTLB	存最近的 IOVA 到物理地址映射的缓存	翻译一次要查多级页表，太慢，所以要有它
RDMA	一台机器直接读写另一台机器内存的协议，两端操作系统都不参与	把它变成攻击工具以后，攻击半径从本机拉到网络
RNIC	支持 RDMA 的网卡	能精确地一次发一条 DMA 请求，逆向 IOTLB 全靠它
mkey	保护 RDMA 内存区域的密钥，随请求一起发送	让对方网卡能验证请求来源
Eviction set	一组地址，访问它们能把目标缓存组里的条目全挤掉	攻击的抓手，和 CPU 缓存侧信道用的是同一套方法
Prime+Probe	先用 eviction set 填满目标组，等一会儿，再访问一次并计时	不需要攻击者与受害者共享内存
Flush+Reload / Evict+Reload	先把目标从缓存里清掉或挤掉，等受害者执行，再读一次并计时	需要双方能碰到同一块内存
DDIO	让设备直接读写 CPU 末级缓存的技术	前人的同类攻击多数打的是它，这篇打的是别处

🧭 起点：让设备来发请求

先交代这篇和 CPU 侧侧信道的根本差别。

CPU 缓存与 TLB 攻击里，挑起事端的是 CPU：攻击者的进程去访问内存，制造缓存冲突。这类攻击有一个共同的前提——攻击者与受害者至少共享某一级存储。而现代云里这个前提常常不成立：每个租户分到独立的 GPU、独立的网卡，通过 PCI-passthrough（把整块物理设备直接划给一台虚拟机）直通进来，内存空间彼此隔离，操作系统层面看不到任何共享。

这篇论文的观察是：即使内存完全隔离，设备之间的地址翻译仍然是共享的。一个 IOMMU 管着一组设备，所有设备的翻译请求都落进同一个 IOTLB。而翻译命中与不命中之间存在 1000 个 CPU 周期量级的时延差，足以测量。

更关键的是攻击者的身份变了。既然请求由设备发出，整个攻击过程不碰 CPU 缓存，也不碰 TLB。论文反复强调这一点，因为它的直接后果是：那些靠 CPU 缓存行为做异常检测的防护，看不见这种攻击。

🔬 先把 IOTLB 拆开看

要攻击一块缓存，得先知道它长什么样。IOTLB 在此之前从没被公开逆向过，所以论文自己动手做了一遍。目标是一颗 Intel Xeon Silver 4316（第三代 Xeon 可扩展处理器）。

观测手段是 Intel 从第三代 Xeon 开始提供的 IOMMU 性能监控（PMON）——CPU 里的一组硬件计数器，可以指定监听哪类事件，同一个事件还能用 umask 这个掩码再细分。其中与 IOTLB 相关的三个事件：

事件	umask	含义
FIRST_LOOKUP	0x1	查 IOTLB 的请求数
MISSES	0x20	未命中后填入新条目的次数
CTX_CACHE_LOOKUP	0x40	未命中后去查上下文缓存（记录每个设备翻译上下文的表）的次数

论文选 MISSES 做判据，因为它的口径更严：只有页表里有有效翻译的未命中才会被计入。

而发起访问的工具选的是 RNIC。理由是它能精确地一次只发一条 DMA 请求，这一点对逐条试探 IOTLB 的组织结构至关重要——GPU 做不到这个粒度。

逆向出来的结果如下：

属性	4 KB 页	2 MB 页
组数（sets）	32	16
路数（ways）	8	8
映射方式	线性（直接拿地址的几位当组号）	线性（同左）
取地址哪几位当组号（中括号表示取第几位到第几位）	IOVA[16:12]	IOVA[24:21]
替换策略	LRU	LRU

这张表里有两条值得单独说。

第一条是大页与小页用的是两套独立的硬件。论文先做了个对照实验：用 4 KB 页把整个 IOTLB 填满，然后拿一个 2 MB 页去挤——一个条目都没被挤掉。这说明两者互不干扰，想让请求之间产生冲突，页大小必须一致。

第二条是替换策略可以被完整确定下来，靠的是前人提出的置换向量方法：准备两个 eviction set，一个用来填满目标组，另一个用来逐条挤；数一数目标条目是在第几次才被挤掉，就能推出它在 LRU 序列里的位置。实验给出的置换向量显示，每个条目的相对位置都被严格维持，IOTLB 用的是完整的 LRU（Least Recently Used，最久没被访问的条目最先被挤出去）。

🎓 先补一课：探测缓存的三种套路

前面说清了 IOTLB 是一块共享缓存，也确认了它的命中与未命中之间有时延差。但有差和能用之间还隔着一步：具体怎么探。这篇论文用的正是业界三种经典套路里的两种，先把它们一次讲明白，后面读到 PP、ER、PPA 这些缩写才不会绊住。

先交代缓存扮演的角色，用图书馆打比方：

主存：图书馆的大书库，书最全，但取一次要走很远；
缓存：你桌子边的小书架，只能放几本，但伸手就能拿到；
命中 / 未命中：要的书在桌上就是命中（快），不在就得去书库翻（慢）。

三种套路建立在同一个前提上：命中快、未命中慢，而这个快慢攻击者能靠计时测出来。区别只在动手前先做什么、最后量什么：

1
2
3

1. Prime+Probe：先把自己的东西塞满书架，过一会儿再看还在不在
2. Flush+Reload：先把某本书从书架上拿走，过一会儿再看它回没回来
3. Evict+Reload：也看书回没回来，但拿走的方式换成用别的书把它挤掉

① Prime+Probe：查自己的地盘有没有被抢

三步：先用自己的数据把目标那一组缓存填满（Prime），隔一段时间让别人干活（Wait），再回来访问自己那批数据并计时（Probe）。如果变慢了，说明自己填进去的条目被挤出去了——而能挤掉它的，只有这期间用过同一组缓存的程序。

它最大的好处是不需要双方共享任何数据：攻击者不用碰到受害者的内存，只要两者的访问会落进同一个缓存组就够。

② Flush+Reload：查某一本书回没回来

三步：先把目标数据对应的缓存行清掉（Flush），等别人干活（Wait），再读一次同一份数据并计时（Reload）。读得快，说明这期间有人把它重新放回了缓存。

它比 Prime+Probe 更精准——盯的是指定的那份数据，而不是笼统的某一组被用过；代价是要求双方共享同一份数据，最常见的情形是共享库里的代码。

③ Evict+Reload：同上，只是不用专门的清除指令

三步和前一种一样，唯一区别在第一步：不用 CPU 提供的清除指令，而是访问一批会争抢同一组缓存的地址，把目标挤出去。好处是不依赖那条特殊指令（有些平台用不了），代价是要自己找出一批能挤掉目标的地址。

三种放在一起

	Prime+Probe	Flush+Reload	Evict+Reload
第一步	填满目标缓存组	清除指定缓存行	靠竞争把目标挤出去
最后量什么	自己那批数据的访问速度	目标数据的访问速度	目标数据的访问速度
需要共享目标数据吗	不需要	通常需要	通常需要
主要看出什么	哪一组缓存被人用过	指定数据有没有被访问	指定数据有没有被访问
特点	适用面广	精度高	不依赖专门的清除指令

回到本文：IOTLB 是设备侧的缓存，上面这三件事照样能做，只是发起访问的从 CPU 换成了 GPU 或网卡。至于这篇具体挑了哪几种、为什么挑它们，下一节说。

如果只记一句：Prime+Probe 看的是自己那批资源有没有被抢走，Flush+Reload 和 Evict+Reload 看的是指定的那份数据有没有被别人放回来。

🛠️ 两个原语：GPU 版与网卡版

有了 IOTLB 的结构参数，攻击就需要一个能持续发起 DMA 请求并计时的抓手。论文做了两套，分别用 GPU 和 RNIC。

GPU 版（DEVIOus_GPU）用 cudaMemcpy() 触发 GPU 到主存的 DMA（CUDA 是 NVIDIA 的 GPU 编程接口，cudaMemcpy 就是它用来在 CPU 与 GPU 之间搬数据的函数）。这里有个细节：CUDA 默认会用临时分配的内存做中转，地址每次都可能变，而攻击需要地址固定，所以必须用 cudaMallocHost() 显式申请固定内存。

而且 GPU 版只能用 Prime+Probe。原因是 GPU 驱动不允许两块内存被不同应用共享，而 Flush+Reload 和 Evict+Reload 都要求攻击者与受害者能碰到同一块内存。

网卡版（DEVIOus_RNIC）走 RDMA，用 ibv_post_send() 和 ibv_poll_cq() 发请求并确认完成。它比 GPU 版多两个好处：一是 RDMA 的内存区域可以在两个客户端之间共享，所以 Evict+Reload 也能用；二是它天然把攻击半径拉到了网络上，后面那两节远程攻击就是用它做的。

🔊 一个技巧：把 1 次未命中放大成 8 次

这部分是全文最精巧的设计，也是远程攻击能成立的关键。

本地测量噪声小，一次未命中带来的延迟差就够用了。远程不一样——隔着 RDMA 网络，时延抖动本身就可能吃掉信号。论文的应对办法不是做统计平滑，而是利用 LRU 的顺序，把信号本身放大。

Prime 之后，一个组里的 8 个条目按 LRU 排好，8 是最老的那个：

  1   2   3   4   5   6   7   8
                              ← 8 在最老的一端，下一个被挤掉的就是它

受害者访问一次新页面 → 8 被挤掉，新条目 T 补进来：

  1   2   3   4   5   6   7   T

反序探测（7 到 1）：只撞上 1 次未命中
正序探测（1 到 7）：每条都比预期慢，撞上 8 次未命中   ← PPA 用这个

命名上，反序探测的叫 Prime+Probe（PP），正序探测的叫 Prime+Probe with Amplification（PPA）。用 PPA 的结果是：探测一轮能撞上 8 次未命中，而不是 1 次。实测里两者的差距很可观——网卡版 PP 的命中与未命中时延差是 1145 个 CPU 周期，PPA 则达到 6323 个周期。

代价也很实在：放大是拿攻击执行时间换来的，而且正是因为它把噪声也一起放大了，PPA 只在噪声大的远程场景里才划算。本地场景里，GPU 版 PP 的 1642 个周期已经比网卡版的任何一个都大，直接用 PP 就好。

原语	方法	命中时延	未命中时延	时延差	单次执行耗时
DEVIOus_GPU	PP	175,633	177,275	1,642	60.42 µs
DEVIOus_RNIC	Evict+Reload	16,798	17,891	983	52.29 µs
DEVIOus_RNIC	PP	147,196	148,341	1,145	45.43 µs
DEVIOus_RNIC	PPA	152,374	158,697	6,323	47.45 µs

（单位都是 CPU 周期，最后一列是单次攻击执行时间。ER 与 PP 的时延绝对值差了一个数量级，是因为两者的请求路径不同，有意义的是差值那一列。）

🎯 四个案例

论文用这两个原语做了四个攻击，从本机跨 VM 一路做到远程网络。

案例一：跨 VM 的隐蔽信道（本地）

两台虚拟机各自直通一张 GPU，共享同一个 IOTLB。发送方用 GPU 的 DMA 时序编码比特，接收方用另一张 GPU 做 Prime+Probe。论文报出的结果是：单个 IOTLB 组、时间窗口 250 µs 时吞吐 2.54 Kbps（千比特每秒）、误码率 2.18%；用上 16 个组以后吞吐升到 40.64 Kbps。

案例二：SSH 击键时序（本地跨 VM）

这个场景值得一提，因为它把攻击角度选得很刁。虚拟机 A 里的攻击者用 GPU 探测；虚拟机 B 里跑的是一台 SSH 服务器，网卡直通给它。攻击者不需要任何 SSH 账号，也不需要访问 B 的权限，只要共享 IOMMU 就够。

原理是利用 Linux 网卡的 RX ring buffer——收包描述符排成的环形队列。网卡每收到一个包，就要按描述符里的地址往主存做一次 DMA，写进接收缓冲区。这个 DMA 会查 IOTLB，于是攻击者能观测到。而且这个 ring buffer 是固定的 4 KB 页、用一致（coherent）方式映射（CPU 和设备看到的是同一份数据，硬件自动保持两边同步），所以活动始终落在同一个 IOTLB 组上——论文用 256 轮发包实验确认了这一点，活跃的永远是 0 号组。

剩下的就顺理成章了：反复探测 0 号组，每个包到达就在时间轴上留下一个峰值。论文的评估是：与抓包工具 tcpdump 的时间戳相比，平均差 0.81 ms，而人类击键的平均间隔约 240 ms——不到千分之四。论文据此认为，这种击键时序攻击已经和嗅探式攻击同等准确。

案例三：远程推断键值访问（RDMA 网络）

攻击者在远端 RDMA 客户端上，目标是服务器上跑着 Apache Crail 的键值存储，想推断另一个客户端访问了哪些键值。这里出现了 PPA 真正的用武之地：

方法	准确率
PP	即使收集 100 次探测记录（trace）也只有 62.66%
PPA	收集 20 条 trace 就从 89.47% 升到 99.18%

同一个噪声环境里，PP 怎么加数据都上不去，PPA 用五分之一的数据就接近满分——放大信号的价值就在这里。

案例四：远程网站指纹（RDMA 网络）

攻击者要推断远端用户在访问哪个网站。抓手是现代浏览器默认打开的 GPU 加速：把网页画成像素这一步（光栅化）和把各层画面拼到一起这一步（合成）现在都交给 GPU，而且一有可用元素就立刻送过去，这个过程产生的 DMA 访问模式与网页的内容和形状相关。

时间上也来得及：加载一个网页平均要 600 ms，而用 PPA 扫一遍整个 IOTLB 只要不到 2 ms，所以一次页面加载能采到 300 多条 trace。分类用的是带注意力机制的 Bi-LSTM（双向长短期记忆网络，一种专治时序数据的模型，论文把它叫作 Att-BLSTM），训练数据取自 Alexa 排名的前 300 个网站。

结果是对 Chrome 和 Firefox、在 Windows 与 Linux 上都做了测试，准确率最高 95.2%；即便换成 Tor 浏览器，仍能到 90.6%。论文另外测了用大页的情况：换成 2 MB 大页时准确率降到 89.52%，作者归因于大页在 IOTLB 里承载的信息粒度更粗。

🛡️ 为什么它比 CPU 侧信道难防

论文在防御部分给了一个很直接的判断：针对 CPU 缓存侧信道的那一整套检测手段，对这套攻击全都无效。

原因在前文已经交代过——攻击请求由设备发出，不经过 CPU 缓存，也不经过 TLB。异常检测器盯着的是 CPU 缓存与 TLB 的行为，自然什么也看不到。

论文给的两条出路：

异常检测要用对了事件。Intel 从第三代 Xeon 起提供了 IOTLB 的硬件事件，论文点名了四个：FIRST_LOOKUPS（umask 0x1）、4K_HITS（0x4）、2M_HITS（0x8）、MISSES（0x20）。用它们来统计异常模式，论文自评的运行时开销不到 5%；
IOTLB 分区。既然风险来自共享，就把 IOTLB 按设备域切开来。论文在另一处（针对同类攻击的讨论里）也承认这条路的代价——隔离会限制规模：一台有 8 个满速插槽、却只有 4 块 IOMMU 的主板，配不出 8 个彼此隔离的单卡实例。

⚠️ 这篇的边界

位置	要注意的地方
前提条件	攻击者设备必须与受害者设备共享同一个 IOMMU。论文的原话是要求同一 root port（PCIe 从 CPU 出来的根端口），比同期一些工作要求共享 PCIe switch（接在 root port 下面、再分叉给多块设备的交换芯片）的假设更宽松
攻击者要有什么	一台能持续发起 DMA 并计时的设备——GPU、RNIC 都可以，或者一块可编程 FPGA。反过来，如果目标应用改用 MMIO（让 CPU 逐次去读写设备寄存器，而不是让设备自己搬数据）而不是 DMA 来驱动设备，这条信道就看不见它
具体数字都绑在平台上	32 组 8 路、LRU、983 到 6323 周期的时延差，全部测自 Intel Xeon Silver 4316 加特定 GPU／网卡组合。换 CPU、换设备都要重新标定
大页会变粗	大页在 IOTLB 里的组数更少（16 组对 32 组），承载的信息粒度更粗——网站指纹的准确率从 95.2% 掉到 89.52% 就是证据
准确率的口径	Crail 那组 99.18% 是在攻击者能反复触发受害者请求的前提下得到的（论文明说这在真实 Crail 里未必可行）；网站指纹用的是闭集设定，即受害者只在 300 个候选网站里浏览

✅ 总结

把发起请求的一方从 CPU 换成设备，IOTLB 就变成了一条不经过 CPU 缓存与 TLB 的共享信道：攻击者不需要与受害者共享任何内存，就能从翻译命中时延的变化里读出隔壁在收包、在访问键值、在渲染网页。

三个可以带走的要点：

共享点的清单需要更新。除了缓存和 TLB，设备侧的地址翻译缓存同样跨租户共享，而且它挡在每一条 DMA 路径上；
换成设备发起请求，就绕开了整类防御。既不碰 CPU 缓存也不碰 TLB，常用的一整套侧信道异常检测在这条路上失效；
噪声大反而催生了一个更聪明的技巧。PPA 不去做平滑，而是利用 LRU 的顺序把一次未命中放大成八次——远程攻击的数据里，PP 加多少数据都停在 62.66%，PPA 用五分之一的数据就到了 99.18%。