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保障团队工作中发现。自2023年9月以来，我们与CERT协调中心合作，进行了大规模协调披露，涉及所有主要GPU供应商，包括NVIDIA、苹果、AMD、Arm、英特尔、高通和Imagination。

截至撰写时，受影响供应商的状态如下：

苹果：尽管多次尝试通过CERT/CC建立联系，我们仅在2024年1月13日收到苹果的回复。重新测试显示部分设备已修补（如Apple iPad Air 3rd G (A12)），但问题仍存在于Apple MacBook Air (M2)。苹果确认A17和M3系列处理器包含修复，但未通知具体补丁部署情况。
AMD：已确认设备仍受影响，正在调查缓解计划。
高通：已发布针对部分设备的固件v2.07补丁，但其他设备可能仍受影响。
Imagination：Google确认部分Imagination GPU受影响，Imagination在2023年12月发布的DDK 23.3中提供了修复。

其他供应商详情：

NVIDIA：确认设备目前未受影响。
ARM：确认设备目前未受影响。
英特尔：测试未观察到受影响。

漏洞简介

GPU最初是为加速图形计算而开发的。在该领域，性能至关重要，先前未发现的安全问题通常对应用程序没有重大影响。历史上，这导致GPU硬件和软件堆栈快速迭代，频繁进行主要架构和编程模型更改，从而产生了复杂的系统堆栈和模糊的规范。

利用要求

这是一种共驻攻击，威胁参与者的攻击途径可以作为共享机器上的另一个应用程序、应用或用户实现。攻击者只需能够运行GPU计算应用程序（例如通过OpenCL、Vulkan或Metal）。这些框架得到良好支持，通常不需要提升权限。攻击者可以通过编写转储未初始化本地内存的GPU内核来读取受害者留在GPU本地内存中的数据。这些攻击程序可以少于10行代码，实现起来并不困难。

除非用户检查应用程序的低级GPU源代码，否则无法发现其应用程序是否使用GPU本地内存；GPU代码通常隐藏在库调用深处，位于深层软件堆栈的低层（例如ML）。总体而言，观察攻击者当前是否正在窃取数据或已窃取数据的方法非常有限。

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

用户缓解措施

鉴于受影响GPU供应商缺乏全面补丁，可以通过修改所有使用本地内存的GPU内核的源代码来防御LeftoverLocals。在内核结束之前，GPU线程应清除内核中使用的任何本地内存位置（例如存储0）。此外，用户应确保编译器不会删除这些内存清除指令（例如通过将本地内存注释为volatile），因为编译器可能会检测到清除的内存未在内核中后续使用。

由于GPU二进制文件通常不显式存储，且GPU二进制分析工具很少，因此难以验证。因此，此缓解措施可能对许多用户来说很困难。

漏洞：LeftoverLocals

我们发现几个GPU框架没有以传统CPU框架预期的方式充分隔离内存。在受影响的GPU上，一个内核（可能来自同一台机器上的另一个用户）可以观察由另一个内核写入的本地内存中的值。因此，通过可编程接口（例如OpenCL）访问共享GPU的攻击者可以从其他用户和进程窃取内存，违反传统的进程隔离属性。

先前学术工作表明，NVIDIA GPU通过各种内存区域（包括本地内存）跨进程泄漏内存。然而，他们仅检查了NVIDIA的GPU，并且未讨论对广泛部署用例（如ML）的影响。LeftoverLocals显示许多GPU仍然容易受到本地内存泄漏的攻击，并且该漏洞可以在重要ML应用程序的共驻攻击中被利用。

总体而言，该漏洞可以使用两个简单程序说明：监听器和写入器。写入器将canary值存储在本地内存中，而监听器读取未初始化的本地内存以检查canary值。

监听器

监听器启动一个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];
  }
}

写入器

写入器启动一个内核，将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];
  }
}

监听器和写入器的CPU程序重复启动各自的内核。如果监听器能够可靠地读取canary值，则该平台易受LeftoverLocals攻击。

窃听LLM响应

恶意参与者（攻击者）可以利用LeftoverLocals窃听多租户GPU机器上另一个用户（受害者）的LLM响应。攻击过程实现上述监听器，并附加步骤将窃取的值与各种指纹进行比较。受害者进程在不知情的情况下成为写入器，写入的值是交互式LLM聊天会话的敏感组件。

攻击最终遵循两个步骤：

攻击过程通过重复转储（即监听）剩余的本地内存来指纹识别受害者进程使用的模型，该内存包含受害者在LLM模型架构中使用的线性代数操作的敏感组件。
攻击者然后重复监听受害者的进程，特别寻求LLM执行输出层，可以通过早期指纹识别的权重或内存布局模式来识别。

输出层是矩阵-向量乘法，有两个输入：模型权重和层输入。由于输出层的模型权重太大无法全面窃取，攻击者可以检查可用的开源模型，通过暴露的模型指纹完全获取权重。我们发现最后一层的第二个输入（即层输入）随后足够小，可以放入本地内存。因此，整个层输入可以被窃取，攻击者可以重现最终层计算以揭示DNN的最终结果。

配置

我们的攻击基于llama.cpp LLM，因其简单性和对GPU加速的各种支持。我们使用易受LeftoverLocals攻击的大型独立GPU：AMD Radeon RX 7900 XT。配置llama.cpp使用OpenCL进行GPU加速，使用CLBLAST线性代数库。使用wizardLM-7B.ggmlv3.q5_0.bin模型，可从Hugging Face获取。

修改

攻击需要矩阵-向量乘法的优化GPU实现。我们发现llama.cpp中当前的矩阵-向量乘法（不调用CLBLAST）未以优化的惯用方式实现。我们将llama.cpp矩阵-向量乘法替换为我们自己的更惯用的实现。

步骤1—指纹识别模型

攻击者如果可以监听受害者的几个推理查询，就可以指纹识别模型。在我们的配置中，GPU包含大约5MB的本地内存。模型有大约33层，每层包含矩阵乘法操作。矩阵乘法通常在GPU上使用平铺进行优化：一种将矩阵细分为小矩阵、执行乘法然后组合结果的方法。在许多优化库中，包括CLBLAST，本地内存用于缓存较小的矩阵。

对于整个推理计算（33层），攻击者可以窃取约80MB的权重，这足以指纹识别模型（假设用户使用开源模型，例如Hugging Face上的模型）。因此，攻击者可以获取受害者使用的完整模型。

步骤2—窃听LLM输出

攻击者然后可以关注DNN的输出层。在我们的配置中，输出层是矩阵-向量乘法，而不是矩阵-矩阵乘法。权重矩阵很大（约128MB），但输入向量很小（约4KB）。由于攻击者在步骤1中指纹识别了模型，攻击者不需要全面窃取权重，因为它们可从指纹识别的模型中获得。

矩阵-向量乘法具有与矩阵-矩阵乘法不同的GPU实现。在输入向量适合本地内存的情况下，最性能的实现通常是将输入向量缓存在本地内存中，因为它被重复使用（即用于重复的点积）。由于输入向量完全存储在本地内存中，攻击者可以窃取整个向量。

在确定攻击者是否找到来自输出层的本地内存时，我们发现攻击者可以简单地查找两侧为零的4KB浮点值。在我们的测试中，这种独特指纹几乎每次都与输出层相关联。对于不同的模型和GPU，此指纹可能需要重新校准。

整合

攻击者拥有权重和输入向量后，可以执行最终计算并获得推理结果。这允许攻击者高保真地重现受害者LLM聊天会话的输出。在实践中，我们调整攻击者非常高效地转储本地内存（即仅使用少量线程并需要少量内存）。这允许攻击者窃听长聊天查询，只有少量明显伪影。

观察到的伪影包括：

重复令牌：由于攻击者进程连续调度两次等情况，攻击者窃取相同的输出层两次，因此LLM未调度计算其下一个令牌。
缺失令牌：攻击者内核未在正确时间调度，即紧接输出层计算内核之后。
输出不正确令牌：由于攻击者错误识别窃取的数据集为最后一层，将打印垃圾令牌；或生成“接近”原始输出的令牌，即使不精确。

尽管存在这些差异伪影，窃取的文本足以揭示LLM响应。此外，攻击者可以通过例如让多个线程启动监听器内核或具有更精确的最后一层指纹来进一步调整。

测试GPU平台的LeftoverLocals

鉴于我们测试设备的多样性，存在几个用各种框架编写的测试LeftoverLocals的应用程序：

Vulkan命令行：使用Vulkan的命令行应用程序。内核用OpenCL编写，并使用clspv编译为SPIR-V。使用名为EasyVK的简单Vulkan包装器。
OpenCL命令行：使用OpenCL框架的命令行应用程序。
苹果应用：可部署在iOS或Mac OS上的苹果应用。使用苹果的Metal框架定位GPU。
Android应用：使用Vulkan定位移动GPU的Android应用。代码使用Vulkan的C API（再次通过EasyVK）使用JNI。内核与Vulkan命令行应用程序相同：用OpenCL编写，并使用clspv编译为SPIR-V。

使用上述程序，我们测试了跨越七个GPU供应商的11台设备（在某些情况下测试多个GPU框架）。我们在三个供应商（苹果、高通和AMD）的设备上观察到LeftoverLocals。泄漏的内存量取决于GPU的大小。较大的GPU包含更多物理内存，因此泄漏更多数据。对于较大的GPU（例如AMD Radeon RX 7900 XT），我们发现每个内核可以泄漏超过5MB。

对于某些设备，特别是Arm的设备，我们未能在监听器中观察到来自写入器的canary值，但确实观察到非零数据。Arm代表审查了我们的观察结果，并得出结论，尽管这些值不为零，但它们不是来自内存泄漏。

此外，我们测试了来自NVIDIA、英特尔和Imagination的一些GPU。对于这些设备，我们仅在本地内存中观察到零，因此未观察到LeftoverLocals。不清楚是否所有设备都未受影响。例如，尽管我们在Imagination设备上未观察到问题，但Google通知我们他们能够在其他Imagination设备上观察到它。

易受攻击的环境

攻击程序必须共驻在同一台机器上，并且必须在受害者在GPU上运行敏感应用程序时同时“监听”。这可能发生在许多场景中：例如，如果攻击程序与受害者共驻在具有GPU的共享云计算机上。在移动设备上，攻击可以在应用或库中实现。监听可以高效地实现，因此可以重复和持续地进行，几乎没有任何明显的性能下降。

接下来，我们简要讨论其他部署GPU或攻击者可能访问敏感信息的环境。尽管某些当前系统（例如WebGPU）目前似乎未受影响，但ML的日益普及和现代GPU的多样性意味着这些系统的下一次迭代（或其他近未来系统）可能会严重受到这些类型漏洞的影响。

云提供商

云提供商（例如AWS和Azure）不太可能提供共享GPU实例，特别是如果用户具有对GPU机器的专用访问权限。在其他情况下，GPU可以使用非常保守的GPU VM技术（例如NVIDIA的vGPU或MxGPU）共享，这些技术物理分区GPU，从而防止用户共享GPU资源（例如本地内存）。因此，许多当前云GPU系统可能目前不易受LeftoverLocals攻击；然而，由于普遍缺乏对这些系统规范和实现的可见性，我们没有确凿证据来确定这一点。我们注意到，我们在多用户Linux服务器以及通过传统多处理的桌面（Windows和Mac）系统上观察到了LeftoverLocals。这包括这些系统上的Docker容器。

移动应用程序

在移动领域的实验和探索中，我们仅在非常特定的实例中能够运行并发GPU进程（来自iOS或Android上的不同应用）。也就是说，我们无法在其他应用（例如受害者）在前台运行时在后台运行GPU进程（例如来自恶意监听器应用）。与我们对云提供商的分析一样，我们无法找到明确详细说明这些约束的文档，因此无法明确声称它们是否易受攻击。然而，如上面的视频所示，当恶意监听器应用与受害者应用并排运行，或者恶意监听器应用从受害者应用快速从后台交换到前台时，LeftoverLocals可以被利用。

远程攻击

我们初步调查了源自网站（例如由远程攻击者托管）的攻击可能性。据我们所知，Web应用程序不具备使用GPU图形框架（例如WebGL）监听本地内存所需的低级功能。我们注意到新的WebGPU框架确实提供了允许网页访问本地内存的低级功能。保守地，WebGPU初始化并执行本地内存（和全局内存）的动态数组边界检查，从而缓解此漏洞。然而，这些检查会导致显著开销。为了进一步测试，我们的代码库包含WebGPU中的简单监听器。正如预期，我们仅在本地内存中观察到零，即使在通过其他框架易受LeftoverLocals攻击的设备上。然而，GPU编译器已知脆弱，不难想象找到可能以某种方式绕过这些检查的编译器错误（特别是使用模糊测试技术）。我们的立场是LocalLeftovers应该在较低级别（例如驱动程序）解决。

GPU供应商如何解决此漏洞

为了防御LocalLeftovers，GPU应该在内核调用之间清除其本地内存。虽然这可能会导致一些性能开销，但我们的实验表明，许多GPU供应商（例如NVIDIA、英特尔）目前似乎提供此功能。甚至似乎为受影响的GPU提供了部分此功能。例如，AMD GPU的Mesa驱动程序在计算内核启动后清除本地内存。然而，这种方法有一个根本缺陷，使其易受LeftoverLocals攻击：此内存擦除是用单独的内核完成的，因此，GPU内核队列可能在计算内核和本地内存擦除之间包含恶意监听器，允许监听器窃取内存。相反，计算内核和本地内存擦除需要原子地发生，即不允许任何其他内核在它们之间交错。否则，用户可以尝试先发制人地防御LeftoverLocals，如下一节所述。

缓解措施

鉴于受影响GPU供应商缺乏全面补丁，可以通过修改所有使用本地内存的GPU内核的源代码来防御LeftoverLocals。正如我们先前所述，在内核结束之前，GPU线程应存储0到内核中使用的任何本地内存位置。鉴于GPU任务通常在内核边界交错，这将防止另一个用户能够读取剩余值。我们注意到此缓解措施可能对许多用户来说很困难，特别是因为GPU代码通常埋在复杂的软件堆栈深处（例如ML）。此外，GPU代码可能是高度优化库的一部分（例如ML线性代数例程）。在这些情况下，很难识别本地内存的使用方式，更难以修改内核以将其清零。可能可以增强编译器以添加此功能，类似于WebGPU处理GPU内存访问的方式（如上所述）。这些缓解措施确实具有应考虑的性能开销。另一种钝化缓解措施涉及简单避免多租户GPU环境。

对LLM和GPU平台的影响

LLM安全

我们的PoC攻击仅检查一个应用程序：交互式开源LLM会话。然而，只要有一点创造力，攻击者很可能可以针对许多GPU应用程序，包括隐私敏感领域中使用的应用程序。我们的动机源于最近开源模型的使用和支持增加，通常伴随着声称其“开放性”通过透明性 inherently 带来安全性和安全性。《自然》杂志最近的一篇文章甚至声称只有开源生成AI模型可以“安全地”彻底改变医疗保健，一个安全关键领域。然而，即使开源模型提供了严格审计和评估的机会（它们尚未被审计和评估），它们的部署仍然依赖于闭源堆栈（即GPU）。正如LeftoverLocals所证明的，开源LLM特别容易受到我们的漏洞攻击，因为我们能够指纹识别这些模型以根据需要获取剩余权重。事实上，我们已经观察到关于在受影响GPU供应商合作部署开源模型的公告，包括Hugging Face与AMD的合作、Lamini在AMD GPU上的部署以及高通和Meta针对边缘设备的合作伙伴关系。

通常，ML的引入带来了传统威胁模型未考虑的新攻击面，这些攻击面可能导致隐式和显式访问数据、模型参数或结果输出，增加了系统的整体攻击面。至关重要的是识别和分类直接影响ML模型的新故障模式类别，以及可能危害ML Ops管道的新威胁，正如我们用LeftoverLocals所证明的那样。我们在以下部分讨论GPU特定的威胁影响。

GPU内存泄漏漏洞LeftoverLocals：窃听LLM响应的新型攻击

研究人员发现名为LeftoverLocals的GPU本地内存泄漏漏洞，影响苹果、高通、AMD等多款GPU，攻击者可跨进程窃取LLM对话内容。漏洞编号CVE-2023-4969，已协调披露并发布PoC攻击代码。