LeftoverLocals：通过泄露的GPU本地内存监听LLM响应

漏洞披露

我们披露了LeftoverLocals漏洞：该漏洞允许从苹果、高通、AMD和Imagination GPU上另一个进程创建的GPU本地内存中恢复数据。LeftoverLocals影响GPU应用程序的整体安全状况，尤其对运行在受影响GPU平台上的LLM和ML模型具有重要意义。

通过恢复本地内存（一个优化的GPU内存区域），我们构建了一个PoC，攻击者可以跨进程或容器边界监听另一个用户的交互式LLM会话（例如llama.cpp），如下图所示：

在AMD Radeon RX 7900 XT上，LeftoverLocals每次GPU调用可泄露约5.5 MB数据，当在llama.cpp上运行7B模型时，每次LLM查询累计达约181 MB。这足以高精度重建LLM响应。该漏洞凸显了ML开发堆栈的许多部分存在未知安全风险，且未经安全专家严格审查。

该漏洞由CVE-2023-4969跟踪。它由Tyler Sorensen作为ML/AI保障团队工作的一部分发现。Tyler Sorensen也是UCSC的助理教授。自2023年9月以来，我们一直与CERT协调中心合作，进行大规模协调披露，涉及所有主要GPU供应商，包括：NVIDIA、苹果、AMD、Arm、英特尔、高通和Imagination。

截至撰写时，受影响供应商苹果、AMD和高通的状态如下：

苹果：尽管多次尝试通过CERT/CC建立联系，我们仅在2024年1月13日收到苹果的回复。我们在1月10日重新测试了漏洞，似乎某些设备已打补丁，即苹果iPad Air第3代（A12）。然而，问题似乎仍然存在于苹果MacBook Air（M2）上。此外，最近发布的苹果iPhone 15似乎不像以前版本那样受影响。苹果已确认A17和M3系列处理器包含修复，但我们未收到关于其设备上部署的具体补丁通知。
AMD：我们已与AMD确认他们的设备仍然受影响，尽管他们继续调查潜在的缓解计划。他们关于此问题的声明可在此处阅读。
高通：我们收到通知，高通固件v2.07有一个补丁，针对某些设备解决了LeftoverLocals。然而，此时可能仍有其他设备受影响。高通代表提供了以下评论：“开发支持强大安全和隐私的技术是高通技术的优先事项。我们赞扬Trail of Bits的AI/ML保障组的Tyler Sorensen博士和Heidy Khlaaf博士使用协调披露实践，并正在向客户提供安全更新。我们鼓励最终用户在设备制造商提供安全更新时应用它们。”
Imagination：尽管我们自己未在测试的Imagination GPU上观察到LeftoverLocals，但谷歌已确认某些Imagination GPU确实受影响。Imagination在2023年12月向客户提供的最新DDK版本23.3中发布了修复。

更多细节在“协调披露”中讨论，测试和受影响的设备列表可在“测试GPU平台的LeftoverLocals”中找到。其他供应商向我们提供了以下细节：

NVIDIA：确认他们的设备目前未受影响。一个原因可能是研究人员先前已探索了NVIDIA GPU上的各种内存泄露，因此他们了解这类问题。
ARM：也确认他们的设备目前未受影响。

尽管我们未收到这些供应商的回复，我们测试了至少一个他们的GPU，并未观察到他们受影响：英特尔。

漏洞简介

GPU最初开发用于加速图形计算。在这个领域，性能至关重要，先前未发现的安全问题通常对应用程序没有重大影响。历史上，这导致GPU硬件和软件堆栈快速迭代，频繁进行主要架构和编程模型更改。这导致了复杂的系统堆栈和模糊的规范。例如，虽然CPU ISA有大量文档，但NVIDIA只提供几个简短的表格。这种模糊的规范导致了 alarming 问题，无论是以前还是现在，正如LeftoverLocals所例示。

利用要求

这是一个共驻利用，意味着威胁参与者的攻击途径可以作为共享机器上的另一个应用程序、应用或用户实现。攻击者只需要能够运行GPU计算应用程序，例如通过OpenCL、Vulkan或Metal。这些框架得到良好支持，通常不需要提升权限。使用这些，攻击者可以通过编写一个转储未初始化本地内存的GPU内核来读取受害者留在GPU本地内存中的数据。如我们的代码所示，这些攻击程序可以少于10行代码。因此，实现这些攻击并不困难，业余程序员也可以访问（至少获取被盗数据）。我们注意到浏览器GPU框架（例如WebGPU）目前似乎未受影响，因为它们向GPU内核插入动态内存检查。

除非用户检查应用程序的低级GPU源代码，否则他们无法发现其应用程序是否使用GPU本地内存；这件事进一步复杂化，因为GPU代码通常隐藏在库调用深处，在深层软件堆栈的低层（例如ML）。总体而言，观察攻击者当前是否在窃取数据或已窃取数据的方法非常有限。此攻击依赖于攻击者读取GPU上未初始化的内存，虽然这在技术上是未定义行为，但目前未动态检查或记录。任何额外的防御都会相当侵入性，例如对GPU内核执行代码分析以检查未定义行为。

我们发布了一个利用此漏洞的PoC，以下部分描述其工作原理。

用户缓解措施

鉴于受影响GPU供应商缺乏全面补丁，LeftoverLocals可以通过修改所有使用本地内存的GPU内核的源代码来防御。在内核结束之前，GPU线程应清除内核中使用的任何本地内存位置（例如存储0）。此外，用户应确保编译器不移除这些内存清除指令（例如通过将本地内存注释为volatile），因为编译器可能检测到清除的内存未在内核后续使用。这难以验证，因为GPU二进制通常未显式存储，且GPU二进制分析工具非常少。由于此类原因，我们注意到此缓解措施可能对许多用户困难，我们在下面的“缓解措施”中进一步讨论。

漏洞：LeftoverLocals

在本节中，我们更详细地描述了名为LeftoverLocals的漏洞及相应利用。然后我们详细介绍了跨各种GPU设备的测试活动，发现AMD、苹果和高通的GPU易受LeftoverLocals影响。对于那些不熟悉GPU架构和术语的人，我们在“背景：GPU如何工作”中提供了更深入的级别设置器。我们还注意到，虽然GPU内存泄露并非新事（下面进一步讨论），但LeftoverLocals展示了比先前发现的漏洞更深的影响和更广的范围。

在高层次上，我们发现几个GPU框架未以传统CPU框架预期的方式充分隔离内存。我们观察到，在受影响的GPU上，一个内核（可能来自同一机器上的另一个用户）可以观察由另一个内核写入的本地内存中的值。因此，通过可编程接口（例如OpenCL）访问共享GPU的攻击者可以从其他用户和进程窃取内存，违反传统进程隔离属性。此数据泄露可能具有严重安全后果，特别是考虑到ML系统的兴起，其中本地内存用于存储模型输入、输出和权重。

先前的学术工作显示，NVIDIA GPU通过各种内存区域（包括本地内存）跨进程泄露内存。然而，他们仅检查了NVIDIA的GPU（本文的结果可能是我们未在NVIDIA GPU上观察到LocalLeftovers的部分原因）。他们也未讨论对广泛部署用例（如ML）的影响。其他工作显示了GPU如何泄露图形数据，以及共驻攻击者可以从另一个进程重建部分视觉信息（参见此处、此处和此处记录的一些示例）。尽管有这些先前工作，LeftoverLocals显示许多GPU仍然易受本地内存泄露影响，并且此漏洞可以在重要ML应用程序的共驻攻击中被利用。

总体而言，此漏洞可以使用两个简单程序说明：一个监听器和一个写入器，其中写入器在本地内存中存储canary值，而监听器读取未初始化的本地内存以检查canary值。监听器重复启动一个从未初始化本地内存读取的GPU内核。写入器重复启动一个将canary值写入本地内存的GPU内核。下面，我们演示每个操作如何执行。

监听器

监听器启动一个从未初始化本地内存读取的GPU内核，并将结果存储在持久主内存区域（即全局内存）中。这可以通过以下OpenCL内核完成：

1
2
3
4
5
6


__kernel void listener(__global volatile int *dump) {
  local volatile int lm[LM_SIZE];
  for (int i = get_local_id(0); i < LM_SIZE; i+= get_local_size(0)) {
    dump[((LM_SIZE * get_group_id(0)) + i)] = lm[i];
  }
}

关键字__kernel表示这是GPU内核函数。我们向函数传递一个全局内存数组dump。内核写入此数组的任何内容稍后可以由CPU读取。我们静态声明一个本地内存数组lm，具有预定义大小LM_SIZE（我们设置为每个测试GPU的本地内存最大大小）。此程序技术上包含未定义行为，因为它从未初始化的本地内存读取。因此，我们使用volatile限定符抑制可能优化掉内存访问的激进编译器优化。事实上，我们的代码包含更多代码模式，以进一步阻止编译器优化掉我们的内存转储。此过程更像是试错过程而非科学。

对于每个循环迭代，调用（线程）从本地内存中的位置读取，并将该位置转储到dump数组中的唯一位置。此代码唯一棘手的部分是指引，因为本地内存跨工作组断开，因此工作组本地ID需要映射到dump中的唯一全局ID。该过程利用内置标识符实现此点，此处记录。在内核结束时，dump包含监听器内核开始执行时存储在本地内存中的每个值。因为dump在全局内存区域中，它可以由CPU主机代码检查以查找canary值。

写入器

另一方面，写入器启动一个将canary值写入本地内存的内核（例如，此工作使用值123）。我们在下面显示OpenCL内核代码示例：

1
2
3
4
5
6


__kernel void writer(__global volatile int *canary) {
  local volatile int lm[LM_SIZE];
  for (uint i = get_local_id(0); i < LM_SIZE; i+=get_local_size(0)) {
    lm[i] = canary[i];
  }
}

此代码与监听器非常相似，除了不是转储本地内存，而是写入一个值。在这种情况下，我们从数组canary写入一个值。我们使用一个额外数组，以便编译器不优化掉内存写入（因为它容易对常数值这样做）。在内核结束时，写入器已用canary值填充所有可用本地内存。

监听器和写入器的CPU程序重复启动各自的内核。在监听器的情况下，每次迭代时，CPU分析在本地内存中观察到的值并检查canary值。在服务器上，这两个程序可以由不同用户或在不同Docker容器中运行。在移动设备上，这些例程可以在不同应用中运行。应用可以交换进出焦点以交替读取和写入。如果监听器可以可靠地读取canary值，那么我们说平台易受LeftoverLocals影响。

以下动画显示监听器和写入器如何交互，以及如果本地内存未清除，监听器如何观察来自写入器的值。

监听LLM响应

在本节中，我们概述了恶意参与者（攻击者）如何利用LeftoverLocals监听多租户GPU机器上另一个用户（受害者）的LLM响应，然后详细描述PoC。

在高层次上，两个参与者都作为共驻进程执行。攻击进程实现上述监听器，附加步骤是将被盗值与各种指纹比较。受害者进程不知不觉是写入器，其中写入的值不是canary值，而是交互式LLM聊天会话的敏感组件。攻击最终遵循两个步骤：

攻击进程通过重复转储（即监听）剩余本地内存来指纹识别受害者进程使用的模型，在这种情况下，由受害者

监听LLM响应：通过泄露的GPU本地内存窃取数据

本文披露了LeftoverLocals漏洞，攻击者可通过GPU本地内存泄露跨进程窃取LLM响应数据，影响苹果、高通、AMD等多款GPU，并提供了技术细节和PoC实现。