LeftoverLocals：通过泄漏的GPU本地内存窃听LLM响应

漏洞披露

我们披露了LeftoverLocals漏洞：该漏洞允许从Apple、Qualcomm、AMD和Imagination 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、Apple、AMD、Arm、Intel、Qualcomm和Imagination。

漏洞利用简介

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

利用要求

这是一个共驻留利用，意味着威胁参与者的攻击途径可以作为共享机器上的另一个应用程序、应用或用户实施。攻击者只需要能够运行GPU计算应用程序，例如通过OpenCL、Vulkan或Metal。这些框架得到良好支持，通常不需要提升权限。

使用这些框架，攻击者可以通过编写转储未初始化本地内存的GPU内核来读取受害者留在GPU本地内存中的数据。如我们的代码所示，这些攻击程序可以少于10行代码。因此，实施这些攻击并不困难，业余程序员也可以访问（至少在获取被盗数据方面）。

我们注意到浏览器GPU框架（例如WebGPU）目前似乎不受影响，因为它们向GPU内核插入动态内存检查。

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

漏洞：LeftoverLocals

在本节中，我们更详细地描述了名为LeftoverLocals的漏洞及相应的利用。然后，我们详细介绍了在各种GPU设备上的测试活动，发现AMD、Apple和Qualcomm的GPU易受LeftoverLocals攻击。

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

这种数据泄漏可能产生严重的安全后果，特别是考虑到ML系统的兴起，其中本地内存用于存储模型输入、输出和权重。

总体而言，该漏洞可以使用两个简单程序说明：监听器（Listener）和写入器（Writer），其中写入器将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];
  }
}

关键字__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聊天会话的敏感组件。攻击最终遵循两个步骤：

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

请注意，输出层是具有两个输入的矩阵向量乘法：模型权重和层输入——换句话说，是从用户输入派生并通过深度神经网络（DNN）早期层传播的值。鉴于输出层的模型权重太大而无法全面窃取，攻击者可以检查可用的开源模型，通过暴露的模型指纹完全获取权重。我们发现最后一层的第二个输入（即层输入）随后足够小，可以放入本地内存。因此，可以窃取整个层输入，攻击者可以重现最终层计算以揭示DNN的最终结果。

我们注意到这是一个相当简单的攻击，通过进一步的创造力和 ingenuity，威胁参与者可能能够构建更复杂和复杂的恶意场景，可能以更严重的方式危害ML应用程序。

我们的配置

我们在下表中概述了我们的配置。我们的攻击基于llama.cpp LLM，因为它简单且支持多种GPU加速。在我们的示例中，我们使用发现易受LeftoverLocals攻击的大型独立GPU：AMD Radeon RX 7900 XT。我们配置llama.cpp使用OpenCL进行GPU加速，它使用CLBLAST线性代数库。我们使用wizardLM-7B.ggmlv3.q5_0.bin模型，可以从Hugging Face获取。选择此模型是因为其合理大小，能够快速原型制作和分析；然而，此攻击可转移到许多不同模型。在我们的威胁模型中，我们假设受害者正在交互式聊天会话中使用LLM。

修改

攻击需要矩阵向量乘法的优化GPU实现。我们发现llama.cpp中当前的矩阵向量乘法（不调用CLBLAST）没有以优化的惯用方式实现。它将部分点积结果存储在本地内存中，然后在最后组合它们。虽然有一种使用线性代数实现我们相同结果的更复杂方法，但为了我们的PoC和演示的简单性，我们用更惯用（遵循最佳GPU编程实践）的自己的矩阵向量乘法替换llama.cpp的矩阵向量乘法。

步骤1—指纹识别模型

如果攻击者可以监听受害者的几个推理查询，它可以指纹识别模型。在我们的配置中，GPU包含大约5MB的本地内存。该模型有大约33层，每层由矩阵乘法操作组成。矩阵乘法通常在GPU上通过平铺优化：一种将矩阵细分为小矩阵、执行乘法然后组合结果的方法（如此处详细说明）。在许多优化库中，包括CLBLAST，本地内存用于缓存较小的矩阵。因此，对于每一层，攻击者可以窃取约2.5MB的权重和约2.5MB的输入。虽然这是大量数据，但我们注意到它不足以重构整个计算。这些层中的许多具有100s MB大的权重和输入。

然而，对于整个推理计算（33层），攻击者可以窃取约80MB的权重，这足以指纹识别模型（假设用户使用开源模型，例如可以在Hugging Face上找到的模型）。鉴于此，我们假设指纹识别模型是一项简单的任务，因此攻击者可以获取受害者使用的完整模型。

步骤2—窃听LLM输出

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

矩阵向量乘法具有与矩阵矩阵乘法不同的GPU实现。在输入向量适合本地内存的情况下，最性能的实现通常是将输入向量缓存在本地内存中，因为它被重复使用（即用于重复点积）。因为输入向量完全存储在本地内存中，攻击者可以窃取整个向量。在确定攻击者是否已找到来自输出层的本地内存时，我们发现攻击者可以简单地查找两侧有零的4KB浮点值。在我们的测试中，这个独特指纹几乎每次都与输出层相关联。对于不同模型和不同GPU，此指纹可能必须重新校准。

整合

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

• 重复令牌：当攻击者由于攻击者进程连续调度两次等情况而两次窃取相同输出层时发生，因此LLM未调度计算其下一个令牌。
• 缺失令牌：当攻击者内核未在正确时间调度时发生，即紧接输出层计算内核之后。
• 输出不正确令牌由于：
  • 攻击者错误识别被盗数据集为最后一层。在这种情况下，它将打印垃圾令牌。
  • 产生与原始输出“接近”的令牌，即使不精确。即，攻击者可能无法窃取目标层的确切令牌嵌入。这导致损坏的令牌嵌入，解码时在语义上（在word2vec意义上）与原始令牌相似。作为示例，在开头提供的GIF中，攻击者提取不正确的单词“Facebook”，在语义上类似于生成文本中的其他命名实体令牌（如“Google”和“Amazon”）。

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

测试GPU平台的LeftoverLocals

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

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

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

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

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

以下YouTube视频演示了在不同平台上使用几个不同应用程序的LocalLeftovers的不同接口和示例——即LLM PoC攻击、隐蔽通信通道和搜索canary值。

易受攻击的环境

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

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

云提供商

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

移动应用

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

远程攻击

我们初步调查了源自网站（例如由远程攻击者托管的网站）的攻击可能性。据我们所知，Web应用程序不具有使用GPU图形框架（如WebGL）监听本地内存所需的低级功能。我们