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

本文披露了LeftoverLocals漏洞，该漏洞影响苹果、高通、AMD和Imagination的GPU，允许攻击者跨进程窃取GPU本地内存数据，特别威胁到在这些GPU平台上运行的LLM和机器学习模型的安全。

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

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

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

图1：LeftoverLocals如何用于对交互式LLM聊天会话实施攻击的示意图

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

图2：LeftoverLocals标志：你的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在其最新的DDK版本23.3中发布了修复，该版本于2023年12月向客户提供。

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

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

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

漏洞简介

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

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

图3：监听器和写入器进程，其中写入器将canary值存储在本地内存中，而监听器读取未初始化的本地内存以检查canary值

窃听LLM响应

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

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

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

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

图4：PoC利用的步骤，攻击者进程可以 uncovering 数据以高保真度窃听另一个用户的交互式LLM会话

我们注意到这是一个相当简单的攻击，通过进一步的创造力和 ingenuity，威胁参与者可能能够构建更复杂和复杂的恶意场景，可能以更严重的方式危害ML应用程序。下面我们详细描述PoC，以及在各种GPU平台上执行的配置和测试，以揭示它们对LeftoverLocals的敏感性。

我们的配置：我们在下表中概述了我们的配置。我们的攻击基于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的输入。虽然这是一个显著的数据量，但我们注意到这不足以重建整个计算。这些层中的许多层有数百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框架的命令行应用程序。
  • 苹果应用：可以部署在iOS或Mac OS上的苹果
comments powered by Disqus
© 2020 - 2025 qife
使用 Hugo 构建
主题 Stack 由 Jimmy 设计
本站总访问量次 本站访客数人次