💻 computer science

Portable Acceleration of Learning With Errors KEMs for Post-Quantum Cryptography

本文展示了一种利用 OpenMP Target offloading 实现的基于容错学习（LWE）的密钥封装机制（KEM）的可移植 GPU 实现，证明了单一源代码库可以在避免供应商锁定（vendor lock-in）的同时，在 NVIDIA 和 AMD 加速器上均实现显著的性能加速和能效提升。

原作者： Tiziana Liberati, Nitin Shukla, Simone Rizzo, Elisabetta Boella, Matteo Barbieri, Gabriella Bettonte, Daniele Gregori, Marco Pedicini

发布于 2026-07-13

📖 1 分钟阅读☕ 轻松阅读

CC BY 4.0

原作者： Tiziana Liberati, Nitin Shukla, Simone Rizzo, Elisabetta Boella, Matteo Barbieri, Gabriella Bettonte, Daniele Gregori, Marco Pedicini

原始论文采用 CC BY 4.0 许可（http://creativecommons.org/licenses/by/4.0/）。 ✨ 这是对下方论文的AI生成解释。它不是由作者撰写或认可的。如需技术准确性，请参阅原始论文。阅读完整免责声明

想象一下，你正试图建造一个超级安全的数字保险库，即使是面对未来的“量子”超级计算机也无法被破解。为此，你需要解决一个巨大的、杂乱的数学难题，叫做“带误差的学习”（Learning With Errors，简称 LWE）。这就像是在试图于一堆数百万个微小且带有噪声的磁铁组成的草堆中寻找一根特定的针，而且当你越深入寻找时，这些磁铁就会晃动得越厉害。

问题在于，解决这些谜题在普通计算机上需要耗费极长的时间。这就像是用一把小勺子一次只能移动一粒沙子，却要搬动一座沙山。作者们提出了一个疑问：“如果我们使用一支庞大且高效的施工队（GPU）来搬运这些沙子呢？”

重大发现：一套代码，两支队伍
通常情况下，如果你想雇佣一支施工队，你必须雇佣特定的团队（比如 NVIDIA 的 CUDA），并编写只有他们能理解的指令。如果你想换另一支团队（比如 AMD），你就必须重写整个说明书。这既昂贵又麻烦。

这篇论文表明，你可以编写一套单一的指令集（使用名为 OpenMP Target 的技术），它既能完美适配 NVIDIA 团队，也能适配 AMD 团队。这就像写一份食谱，无论是在燃气灶还是电炉灶的厨房里烹饪，都不需要改变任何配料，食谱依然适用。

神奇的技巧：留在厨房里
这些数学谜题中最浪费时间的部分，是在主计算机（CPU）和超快速 GPU 之间来回奔波。这就像厨师为了每撒一丁点调料，都要跑去储藏室拿一次盐。

作者们想出了办法，让“盐”（数学运算所需的随机数）直接留在 GPU 的厨房里。他们升级了一个名为 RngonGPU 的工具，使其能同时适用于 NVIDIA 和 AMD 团队。现在，GPU 可以生成其所需的所有随机数，而无需离开座位。这让工作流程变得流畅且快速。

竞赛：谁会获胜？
团队在四台不同的超快速计算机上测试了他们的新方法：

NVIDIA A100： 一台强大且标准的超级计算机。
NVIDIA GH200： 一台巨大的“超级芯片”，其大脑（CPU）和肌肉（GPU）通过一条超高速公路（NVLink）粘合在一起。
AMD MI300X： 一台拥有巨大内存库的强大标准超级计算机。
AMD MI300A： 一种大脑和肌肉共享完全相同内存库的芯片。

以下是他们的发现：

速度提升： 当数学谜题变得巨大（规模为 4,096）时，GPU 版本比普通计算机版本快了 120 倍。对于规模更大的谜题（规模为 16,384），GPU 依然运行得飞快，而普通计算机则几乎处于停滞状态。
赢家： NVIDIA GH200 是最快的，大约在 60 秒内完成了任务。AMD MI300X 紧随其后，耗时 85 秒。这两台机器都拥有巨大的、高速的内存库（HBM3），可以快速倾泻数据。
令人意外的输家： 你可能会认为 AMD MI300A 会表现出色，因为它实现了大脑和肌肉共享内存。但它实际上是最慢的，耗时 114 秒。为什么呢？因为大脑和肌肉在争夺同一根水管。当 GPU 试图移动数据时，CPU 也在进行自己的数学运算，它们互相干扰。这就像两个人同时试图用同一根吸管喝水；谁也无法喝得痛快。

能量与热量
团队还检查了这些机器消耗了多少电量。NVIDIA GH200 不仅完成得更快，而且消耗的能量更少。它用了大约 9.7 kJ 的能量来解决谜题，而 AMD MI300X 需要约 26.2 kJ。这意味着 NVIDIA 的机器在能量效率上大约高出 2.5 倍。

这意味着什么
这篇论文证明了，你不需要为了获得超高速安全性能而在 NVIDIA 和 AMD 之间做单选题。你可以使用一套代码在两者之上运行。然而，他们也表明，仅仅拥有快速的芯片是不够的；内存是如何组织的同样重要。如果大脑和肌肉在争夺同一块内存，整个系统就会变慢。

简而言之，作者们衡量出，可移植的、经 GPU 加速的安全技术不仅是一个梦想，它已经成为现实，能够让量子防范锁变得更快、更高效——只要你选择正确的硬件配置，以避免在内存通道上发生交通拥堵。

技术摘要：面向后量子密码学的可移植学习误差（LWE）密钥封装机制（KEM）加速研究

问题陈述
向后量子密码学（PQC）的过渡要求实现能够满足现实应用计算需求的方案。在 PQC 候选算法中，基于学习误差（LWE）的密钥封装机制（KEM）提供了强大的安全基础，但会产生巨大的计算开销。这种开销主要由大规模矩阵-向量运算以及加密安全随机数的生成所驱动。虽然图形处理器（GPU）是加速这些并行工作负载的天然选择，但现有的 GPU 实现很大程度上依赖于 CUDA 特有的优化，这造成了厂商锁定，并限制了在异构加速器生态系统（如 NVIDIA 和 AMD）中的可移植性。此外，以往使用 OpenACC 的可移植尝试显示出跨厂商支持有限的问题。本文旨在解决的核心挑战是：如何在保持关键组件（如随机数生成）驻留在设备端的同时，实现高性能的稠密 LWE 型 KEM 加速，并维持源代码级的跨不同 GPU 架构的可移植性。

方法论
作者展示了一种使用 OpenMP Target offloading 实现的纯 LWE 型 KEM 的可移植 GPU 实现。该方法涉及三个主要的策略技术：

可移植卸载策略： 实现利用 OpenMP Target 指令来管理数据移动和内核执行。为了确保 GPU 驻留并最小化主机-设备间的通信，作者采用了持久化设备内存管理（target enter data map(alloc:...)）。这使得主要的矩阵、密文和临时缓冲区在密钥生成（KeyGen）、封装（Encapsulation）和解封装（Decapsulation）阶段都能全程驻留在 GPU 上。
异构随机数生成（RNG）： 工作流中的一个关键组成部分是加密安全的随机数生成。作者扩展了原先为 CUDA 设计的现有 RNGonGPU 库，使其支持异构计算接口（HIP）。通过将 HIP 兼容的适配层集成到 OpenMP Target 工作流中，他们实现了在 NVIDIA 和 AMD 加速器上完全由 GPU 驻留的随机数生成。这是通过使用 #pragma omp target data use_device_ptr 将 OpenMP 管理的设备缓冲区暴露给外部基于 HIP 的 RNG 后端来实现的，从而避免了在主机端生成随机数并将其拷贝至设备端的需要。
内核并行化： 计算内核使用 OpenMP teams distribute parallel for 结构进行并行化处理。密钥生成涉及并行生成秘密向量和误差向量，随后进行稠密矩阵-向量乘法。加密和解密过程采用批处理方式，同时处理独立的比特位或密文槽，以最大化线程级并行度和缓冲区复用。Fujisaki-Okamoto (FO) 变换中的等值检查也在 GPU 上完全执行，仅向主机返回一个状态标志。

实验设置
研究在四种不同的加速器架构上进行了评估：

NVIDIA: A100-SXM-64GB (PCIe 4.0, HBM2e) 和 GH200 Grace Hopper (NVLink-C2C, HBM3)。
AMD: MI300X (PCIe 5.0, HBM3) 和 MI300A (统一 APU, 共享 HBM3)。

性能针对多核 OpenMP CPU 基准进行了测量，并在不同平台间进行了比较，测试的格维度（ $n$ ）最高达 16,384。评估指标包括执行时间、可扩展性、解题能效比（energy-to-solution），以及使用厂商工具（Nsight Systems 和 ROCm Systems Profiler）进行的详细运行时剖析。

关键结果

相对于 CPU 的加速： OpenMP Target offloading 展示了相对于多核 CPU 基准的显著加速。对于大格维度（如 $n=4096$ ），加速比超过两个数量级（120倍），这主要是因为该工作负载属于内存受限型，而 GPU 为流式矩阵运算提供了卓越的内存带宽和并行性。
跨平台性能：
- NVIDIA GH200 取得了最佳整体性能，这归功于其高 HBM3 带宽（4 TB/s）和 NVLink-C2C 互连，后者最小化了 CPU-GPU 通信开销。
- AMD MI300X 紧随其后，在较小的格维度下表现出近乎最优的性能和强大的可扩展性，但在较大的 $n$ 值下，其执行时间略高于 GH200。
- NVIDIA A100 相比 GH200 出现了适度的减速，受限于其较低的 HBM2e 带宽和 PCIe 通信开销。
- AMD MI300A 表现出最弱的可扩展性和最高的执行时间。剖析显示，尽管拥有高带宽的统一内存，但共享的 HBM3 子系统导致了 CPU 端操作（SHA3 哈希、FO 密钥派生）与 GPU 内核之间的竞争，使有效带宽比离散型的 MI300X 降低了约 38%。
能效： 研究发现，对于大格维度，NVIDIA GH200 的能效比 AMD MI300X 高出约 2.5 倍，这由更短的执行时间和更低的平均功耗共同驱动。
可移植性： 单一源代码库成功在所有测试平台上运行，无需针对特定架构修改代码，验证了 OpenMP Target 用于可移植 PQC 加速的可行性。

意义与主张
本文主张，可移植 GPU 加速是减少 PQC 计算开销且不引入厂商锁定的可行策略。研究表明：

源代码级可移植性是可以实现的： OpenMP Target 结合集成了 HIP 的 RNG 后端，允许单一代码库有效地利用异构 GPU 生态系统。
内存架构决定性能： 对于内存受限的 LWE 工作负载，峰值计算能力不如内存带宽、内存系统组织以及主机-设备互连效率重要。结果强调，统一内存架构（如 MI300A）并不自动保证优越性能，如果 CPU 和 GPU 流量争夺相同的内存资源。
实际部署： 通过避免对 CUDA 特有依赖，这种方法促进了在多样化硬件环境中部署量子抗性密码学基础设施。

作者总结道，虽然 OpenMP Target 提供了有效的源代码级可移植性，但运行时性能仍取决于编译器工具链的成熟度、交互式设备库的行为（特别是关于 HIP RNG 适配器中的动态内存分配），以及目标硬件的具体内存层次结构。

技术摘要：面向后量子密码学的可移植学习误差（LWE）密钥封装机制（KEM）加速研究

类似论文