VIPIR：一个用于集成私有信息检索协议的通用GPU框架

大家读完觉得有帮助记得关注和点赞！！！

摘要
私有信息检索（PIR）通过完全隐藏访问模式来实现私有数据库服务，但同时需要高计算吞吐量、大内存容量和充足的内存带宽。我们推出了VIPIR，一个通用的GPU框架，协同设计PIR协议与GPU加速。我们开发了一个统一的分析模型，表明最先进的PIR协议分为两类且具有互补的局限性，并提出了两种能够灵活结合这些类别技术的协议，克服了两类协议的限制。这些协议采用了一种称为基于扩展的环打包（ExpPack）的GPU友好型数据压缩方法，该方法提供了高度的并行性和最小的通信成本。VIPIR对核心操作进行了进一步优化，包括数论变换（NTT）和各种矩阵乘法（GEMM）。值得注意的是，我们开发了一种基于张量核心的数据库乘法执行方法，将其解释为混合整数类型的GEMM。我们还设计了内存高效的调度方法，最大限度地减少了中间缓冲区，并支持在内存容量限制下的多GPU扩展。总体而言，VIPIR实现了比先前PIR系统高数个数量级的吞吐量，同时降低了通信和内存开销，使大规模PIR变得实用。

1. 引言

随着数据隐私成为数字服务的关键挑战，私有信息检索（PIR）[1, 2]已成为一种关键的密码学原语。PIR使客户端能够从远程服务器上的数据库中检索记录，而无需透露正在访问哪条记录。这种能力在实践中越来越重要：例如，以太坊基金会正在探索使用PIR来实现区块链访问的完全匿名化[3]，而苹果公司则将PIR用于保护隐私的联系人发现和安全浏览[4]。

尽管PIR具有理论上的吸引力，但由于一个基本要求，大规模部署PIR仍然具有挑战性：为了保护隐私，服务器必须为每个查询对整个数据库进行线性扫描。单服务器PIR协议[2]在此扫描过程中对数据库执行同态加密（HE）操作[5]；HE允许对加密数据直接进行计算，从而能够在不泄露访问索引的情况下私有地检索记录。这种基于HE的线性扫描对系统提出了严格的要求，同时需要高计算吞吐量、大内存容量和充足的内存带宽。

GPU提供了加速此类扫描所需的大规模并行性。与大多数先前密码学PIR研究所采用的主要平台——服务器级CPU系统相比，现代GPU通过高带宽存储器（HBM）DRAM提供显著更高的内存带宽，并在更低的每操作能耗下提供数个数量级的计算吞吐量。

图1：先前基于CPU的PIR研究和VIPIR针对16GiB/64GiB数据库（单批次）的服务器端在线计算时间和客户端-服务器通信量。VIPIR支持多种协议（混合、双打包、OnionPIRv2、ExpPack（S/D））。InsPIRe和我们的混合协议跨越多种数据库维度。更多细节见表1。

然而，关于PIR的GPU聚焦研究很少[6, 7, 8]，而且它们专注于单一协议，很少考虑使协议设计适应高效的GPU执行。尽管与可配置TB级内存的CPU系统相比，GPU的DRAM容量相对有限，但PIRonGPU [7]和IVE的GPU实现[8]采用的协议会将数据库大小扩展约4倍。这种数据库膨胀成为可扩展性的关键瓶颈。

为了寻求一种有原则的方法，我们分析了现有的PIR协议，并开发了一个统一的分析模型，使我们能够从共同的角度对它们进行推理。尽管有各种各样的方案，但我们观察到，大多数PIR协议在此模型下分为两类：一类（我们称之为标量HE方法）具有较高的客户端-服务器通信开销，另一类（多项式HE方法）存在数据库膨胀问题。

在本文中，我们推出了VIPIR，一个通用的GPU框架，能够实现GPU感知的PIR协议设计，并在NVIDIA GPU上进行优化执行。除了在GPU上加速现有的PIR协议外，我们还从我们的统一模型中推导出新的协议，这些协议本质上更适合高效的GPU执行。这些协议以一种称为环打包（或简称打包）[9, 10, 11, 12]的HE原语为中心，该原语将加密数据压缩到明显更小的尺寸（在我们的设置中为0.16%），以计算换通信。虽然打包已在先前的PIR工作中使用[13, 14, 15, 16]，但我们发现了它在我们的协议中的不同用法，并且还开发了一种新颖的基于扩展的打包（ExpPack）方法。

具体而言，我们引入了（i）一种双打包设计，应用两次环打包以显著提高压缩率，以及（ii）一种混合设计，重新利用环打包在标量HE和多项式HE方法之间进行转换。这避免了多项式HE的数据库膨胀问题，同时减轻了标量HE的通信开销。

ExpPack通过用服务器端扩展的单个密文（96KiB）替换传输n=1,280个“特殊”密文（120MiB），利用GPU的高计算能力以计算换通信，从而增强了这些协议。虽然先前的研究也提出了通信高效的打包方法，但它们依赖于顺序过程或内存密集型预计算，限制了GPU的效率[16, 14]。相比之下，ExpPack旨在通过数千个独立的子操作（即数论变换（NTT）和模矩阵乘法（GEMM））暴露大规模数据并行性，从而有效利用数百个流式多处理器。

VIPIR结合了各种架构感知的GPU优化，以实现实用的性能。值得注意的是，我们利用NVIDIA张量核心[17]通过将INT8-INT32 GEMM模拟到INT8 GEMM数据路径上来加速核心的INT8-INT32 GEMM，这是通过将参数和协议与GPU实现协同设计实现的。我们还调整了数据布局以实现连续内存访问，应用了诸如惰性归约[18, 19]之类的数值技术来减少指令数，并通过在受限GPU内存下仔细安排作业来最小化临时数据。VIPIR还通过结合数据库分片和批量作业分配来支持多GPU执行，从而实现对更大数据库的低延迟扩展。

VIPIR在支持单个GPU（H100 NVL 94GB）上高达64GiB的数据库大小的同时，提供了显著的性能提升，同时避免了由于内存使用过多而导致的先前工作中常见的内存不足故障。如图1所示，VIPIR实现了比先前工作低数个数量级的延迟，我们协同设计的双打包和混合协议同时降低了通信成本。对于一个64GiB的数据库，这些协议实现了33.0-33.5毫秒的执行时间，仅比从DRAM流式传输一次数据库的19.7毫秒硬延迟界限高1.67-1.69倍，表明我们的实现非常接近理论性能极限。VIPIR还实现了163.1-232.8 QPS的高吞吐量，比先前基于GPU的工作ShiftPIR [6]提高了652.6-931.3倍。

我们将VIPIR的主要贡献总结如下：

• 我们开发了一个统一的分析模型，该模型捕获了多种PIR协议，暴露了它们的局限性，并实现了新的双打包和混合设计。
• 我们引入了ExpPack，一种GPU友好的基于扩展的环打包方法，在最小化客户端-服务器通信的同时暴露了大规模并行性。
• 我们构建了优化的PIR GPU实现，包括张量核心加速、内存感知布局和调度、数值优化以及多GPU扩展。

2. 动机与背景

向量（例如，v）和矩阵（例如，A）用粗体表示；所有向量均为列向量。一个 M×K 乘以 K×N 的矩阵乘法表示为 M×N×K GEMM。主要符号和标记总结在附录A中。

2.1 私有信息检索（PIR）的挑战

数十种PIR协议[20, 13, 21, 22, 23, 24, 25, 16, 14, 15, 26, 27, 28]在客户端/服务器端计算、内存访问、存储开销和客户端-服务器通信之间表现出复杂的权衡。尽管如此，大多数构建都基于先前的工作，重用早期协议的核心组件和设计原则。

在没有大量客户端预处理[29]的情况下，PIR协议要求服务器为每个查询线性扫描整个数据库，以隐藏被访问的记录。由此产生的对GB级数据的统一访问需求对计算系统提出了同时提供高带宽和大容量的挑战，而内存层次结构在缓解这一瓶颈方面提供的帮助有限。

在本文中，我们提炼PIR的共同因素，以构建一个有原则的PIR框架，重新思考现代加速器（如GPU）如何重塑设计空间。为了聚焦分析，我们排除了依赖限制性假设的PIR协议，例如多服务器非共谋[1]、单客户端批量查询[30, 21]或大量的客户端数据库预处理[29]，这些假设阻碍了实际应用。

图2：(a) SimplePIR [31] 的服务器端 B 批次执行和 (b) 基于 Tiptoe [15] 的环打包。每个单元格代表一个 INT8/INT32 元素。红框单元格标记客户端交换的数据，其中虚线单元格在打包下不交换。

2.2 同态加密（HE）

HE允许对加密数据直接进行计算，这对于实现PIR至关重要。对于分别加密明文 pt₁ 和 pt₂ 的 HE 密文 ct₁ 和 ct₂，HE方案具有以下线性性质，允许密文间加法和明文-密文乘法：

我们将各种HE方案分为两类：

• 标量HE：将一个整数（明文）加密成一个长度为 (n+1) 的整数向量（密文）。
• 多项式HE：将一个 N 系数整数多项式（明文）加密成环 ℛ_q = ℤ_q[X]/(X^N+1) 中的两个或多个多项式（密文）。

具体来说，标量HE指的是基于带误差学习（LWE）问题[32]的基本Regev加密。我们使用的参数设置是：可以将一个8位整数（INT8）加密成一个32位整数（INT32）向量。多项式HE包括基于环LWE（RLWE）[33]、环GSW（RGSW）[34]或模LWE（MLWE）[35]的各种HE方案。对于最常用的RLWE，每个密文由 ℛ_q 中的两个多项式组成；即 ct ∈ ℛ_q²。我们在整个工作中使用 n=1,280，N=2¹²，q≃2⁸⁷（第5节详述）。

2.3 标量HE算术与多项式HE算术

标量HE的计算相对简单，不需要模算术[36]。对于我们的典型参数，明文-密文乘法和密文间加法分别转换为常规的INT8-INT32标量-向量乘法和INT32向量加法。

处理多项式HE的多项式在计算上要求更高。ℛ_q 中的一个多项式在 ℤ_q 中有 N 个系数，其中使用剩余数系统（RNS）作为优化，将模数 q 分解为 #RNS 个小素数。例如，我们设置 q = q₀·q₁·q₂（#RNS=3），其中 q_i < 2³²，将一个多项式转换为一个 3×N 的INT32矩阵，其中每行对应不同的 q_i：

虽然多项式HE的密文间加法成为简单的逐点模 q_i 矩阵加法，但明文-密文乘法涉及多项式乘法。

数论变换（NTT）是加速多项式乘法的标准技术。对于 3×N 矩阵，NTT对每一行应用傅里叶变换的一种变体（计算复杂度为 O(N log N)）。一旦转换完成，一个多项式乘法（O(N²)）就简化为逐元素模 q_i 乘法（O(N)）。

2.4 密文膨胀与明文膨胀

每个HE类别都面临着一个独特的数据大小爆炸问题：标量HE的密文膨胀和多项式HE的明文膨胀。标量HE加密导致巨大的数据扩展，为 (n+1)·INT32/INT8 = 5,124 倍，我们称之为密文膨胀。相比之下，多项式HE可以利用明文多项式（∈ ℛ_p，本文中 p=2¹⁸）中的 N 个系数，每个密文加密 N log₂ p 比特的数据。特别是，RLWE密文（∈ ℛ_q²）仅产生 2·(log₂ q)/(log₂ p) 倍的扩展，低于10倍。

相反，在多项式HE中，明文在应用NTT以实现高效的明文-密文乘法时会增长。在此过程中，ℛ_p 中的明文被提升到 ℛ_q，导致大小扩展为 (log₂ q)/(log₂ p) 倍，通常约为4倍。这种明文膨胀问题在标量HE中不会出现，其中明文（标量）保持其原始形式。

2.5 示例PIR协议：SimplePIR

作为一个具体例子，我们来看SimplePIR [31]，它将数据库组织成一个 D₁×D₀ 的矩阵。为了查询沿 D₀ 维度的索引 i，客户端发送 D₀ 个标量HE密文，其中只有第 i 个密文加密1，其余加密0，形成 i* 的加密独热表示。这些密文一起形成一个 D₀×(n+1) 的矩阵，服务器将其与数据库相乘，产生一个 D₁×(n+1) 的矩阵，该矩阵被解释为 D₁ 个标量HE密文，加密了数据库的第 i* 列。

SimplePIR利用了 D₀×n 个输入元素与查询无关且可以在客户端之间重用的事实。因此，大部分数据库GEMM可以预计算，产生一个 D₁×n 的提示。如果客户端系统在离线通信阶段接收并存储提示，那么客户端在在线阶段只需发送 D₀ 个元素，并接收 D₁ 个元素作为响应。

图2(a)说明了具有多客户端批处理[8]的SimplePIR的服务器端计算。通过提示预计算，服务器的在线计算仅包括用于 B 个查询批次的 D₁×B×D₀ INT8-INT32 GEMM。

2.6 用于无提示SimplePIR的环打包

许多研究[13, 14, 15, 16]通过应用环打包（或简称打包）探索了SimplePIR的无提示变体，该打包将 N 个具有 N×(n+1) 个元素的标量HE密文转换为一个具有 2·(#RNS·N) 个元素的多项式HE（RLWE）密文。这表示压缩因子（越小越好）为 2·#RNS/(n+1) 倍，减轻了标量HE中的密文膨胀。打包后，#RNS可以进一步减少到1（称为模归约[37]），实现 2/(n+1) 倍（对于我们的参数 <0.16%）的压缩。这消除了对提示重用的需求，而是通过紧凑的多项式HE密文实现在线通信。

图2(b)展示了基于Tiptoe [15]的示例打包，涉及NTT和带有额外“特殊”多项式HE密文的并行GEMM。环打包将在第3.5节中详细分析。

2.7 GPU执行模型

现代NVIDIA GPU是由多个流式多处理器（SM）组成的大规模并行系统，每个SM集成了执行单元，如用于整数算术的INT32核心和用于高吞吐量低比特（例如INT8）GEMM的张量核心。GPU内核遵循单指令多线程（SIMT）模型，其中线程被分组为warp（32个线程一组）并调度到SM。性能由占用率（即每个SM相对于硬件限制的活动warp数量）以及内存层次结构的有效利用决定，内存层次结构涵盖每个SM的内存资源（寄存器、共享内存和L1缓存）和全局内存资源（L2缓存和DRAM）。

最近的GPU采用高带宽存储器（HBM）来提供巨大的DRAM带宽（例如，H100 NVL 94GB中为3.94TB/s），并为多GPU配置提供高带宽互连（如NVLink），从而能够以较小的通信开销实现高效并行化。

3. PIR协议的统一方法

3.1 将PIR视为一个压缩过程

我们通过将PIR解释为一个压缩过程，提出了一个独特的视角。一个平凡的PIR协议让服务器返回整个数据库作为对每个查询的响应，而实用的协议则通过HE操作压缩响应，只保留请求的记录。例如，SimplePIR（第2.5节）实现了 4(n+1)/D₀ 倍的压缩因子，将数据库（D₁×D₀ 个INT8值）压缩成一个紧凑的响应（D₁×(n+1) 个INT32值，包括提示）。

更一般地，PIR协议应用一系列维度折叠（或简称折叠）步骤来逐步压缩数据库。对于一个具有 D_k × D_{k-1} × … × D₀ 条记录的数据库，每个折叠步骤会从响应中去除一个 D_i 因子。

3.2 PIR的统一分析模型

我们观察到PIR协议中的二分法，大多数PIR协议依赖单一HE类别进行维度折叠。关于各个协议的更多细节在第7节提供。

• 标量HE折叠。 基于SimplePIR，这些协议有一个或两个使用独热表示的折叠步骤（第2.5节）；两个折叠的变体称为DoublePIR [31]。第二次折叠需要事先进行重新解释，因为标量HE不支持密文间乘法，将现有数据视为标量（明文）流。无提示变体在最后添加一个环打包步骤（第2.6节）。此PIR类别包括SimplePIR/DoublePIR [31]、Tiptoe [15]、HintlessPIR [14]和YPIR [13]。

• 多项式HE折叠。 此类别的优化协议OnionPIRv2 [28]将数据库组织成 2×2×…×2×D₀ 的结构，利用RLWE明文-密文乘法进行初始 D₀ 维度折叠，并使用RLWE-RGSW外积进行其余折叠[25]。尽管也使用了其他方法，如RLWE密文间乘法[21]、基于旋转的加密矩阵-向量乘法[20]或MLWE操作[24]，但它们提供了类似的计算-通信权衡。此类别包括XPIR [23]、MulPIR [38]、WhisPIR [27]、OnionPIR [25]、OnionPIRv2 [28]、Spiral [24]、Respire [22]和KsPIR [20]。

图3：我们的统一PIR模型。每个箭头标注了其典型的压缩因子（越低越好），表示直接转换响应状态的操作；确切的压缩因子可能因参数和实现而异。这里，D_i 表示数据库维度，n=1,280 是LWE次数。

我们开发了一个克服二分法并提供PIR协议综合视角的统一分析模型（见图3）。在此模型中，两个PIR类别遵循以下路径（可选路径用括号括起）：
• 标量HE折叠：① → (② → ①) → (③)
• 多项式HE折叠：⑤ → ⑥ → ⑦ → ⑦ → … → ⑦

也有一些例外；例如，SealPIR [21]遵循 ⑤ → ⑥ → ④ → ⑤ → ⑥，尽管与其他的多项式HE折叠方法相比，这种设计通常效率较低[25]。此外，InsPIRe [16]开发的协议遵循的路径与我们第3.4节中的新协议类似，但InsPIRe在针对GPU时显示出显著的局限性（第3.5.3节）。

3.3 两个PIR类别的局限性

在此模型下，我们识别出每种方法的局限性。标量HE折叠存在响应大小大的问题。对于 ① → ② → ① → ③，总体压缩因子为：
4(n+1)/D₀ × 1 × 4(n+1)/D₁ × 2/(n+1) = 32(n+1)/(D₀ D₁)。
尽管常数项随参数选择而变化，但源自标量HE中密文膨胀的 (n+1)（≥2¹⁰）项导致相对于原始记录大小的大幅响应扩展。

相比之下，多项式HE折叠实现了更有利的压缩因子（≃ 8/∏ D_i），但遭受数据库膨胀之苦，即由于初始NTT（⑤）导致的明文扩展。由此产生的约4倍的数据库大小增长严重限制了PIR协议对大型数据库的适用性和可扩展性。

数据库膨胀问题对GPU尤其关键，其内存容量（每设备高达数百GB）比服务器级CPU（高达TB级）更有限。由于PIR需要对整个数据库进行线性扫描，将数据库保存在GPU DRAM中对于最大化带宽至关重要。然而，对于现实世界GB到TB规模的数据库，额外的约4倍存储开销成为主要瓶颈，要么导致昂贵的主机-GPU通信，要么仅仅为了内存容量就需要额外的GPU。

图4：基于扩展的环打包（ExpPack）的黑盒视图，用于统一模型中的多条路径；红框表示客户端发送的数据。

3.4 GPU友好的协议构建

从我们的统一模型出发，我们推导出两种PIR协议，它们实现了改进的压缩，降低了客户端-服务器通信，同时通过从标量HE折叠（①）开始，避免了数据库膨胀——这对GPU至关重要：

• 双打包： ① → ③ → ④ → ① → ③。应用两次打包将总体压缩因子降低到 64/(D₀ D₁)，消除了 (n+1) 因子。这实质上是用来减少的通信换取执行两次打包的额外计算。我们通过一种新颖的环打包方法（ExpPack，第3.5节）来减轻计算开销，该方法对多次打包计算特别有效。

• 混合： ① → ③ → ⑦ → ⑦ → … → ⑦。虽然先前的标量HE折叠方法主要使用打包进行压缩，但我们将其重新用作桥梁，以实现混合PIR协议，该协议从初始维度的标量HE折叠过渡到后续维度的多项式HE折叠。对于多项式HE折叠（⑦），我们采用RGSW-RLWE外积[39]，如OnionPIR中所用[25, 28]。

两种协议的总体计算成本主要由前两个步骤（① → ③）决定；在此之后，数据被充分压缩（8/D_i 倍），后续计算贡献甚微。

3.5 基于扩展的环打包（ExpPack）

我们的协议特别适合GPU，部分原因是一种新颖的基于扩展的打包方法（③），该方法在暴露大规模GPU并行性的同时，极大地减少了客户端-服务器通信。在现有的打包方法中[12, 11, 9, 10]，核心思想是相同的：使用客户端提供的额外“特殊”密文（见图2(b)）和/或密钥[16, 12, 14]，将 N 个标量HE密文转换为单个多项式HE（RLWE）密文。

ExpPack的主要区别在于特殊密文的构造方式，使服务器能够即时扩展它们。扩展是一种HE算法，能够从加密了 a₀ + a₁X + … + a_{N-1}X^{N-1} 的RLWE密文中提取出每个加密 a_i（0 ≤ i < N）的RLWE密文[21, 25]，我们在本工作中以黑盒方式使用它。

我们建议将这种扩展与Tiptoe [15]中的打包方法相结合，否则由于特殊密文过高的通信开销而不实用。对于客户端的长度-n 秘密 s = (s₀, s₁, …, s_{n-1})，Tiptoe需要 n 个特殊的RLWE密文用于环打包：
Encrypt(s₀), Encrypt(s₁), …, Encrypt(s_{n-1})。
对于我们的参数，它们总大小达到120MiB，比实际的索引和响应密文（KB级）大几个数量级，使得Tiptoe非常不实用。

相反，我们让客户端发送一个加密了所有秘密值的单个RLWE密文（96KiB）：
ct_ExpPack ← Encrypt(s₀ + s₁X + … + s_{n-1}X^{n-1}), n < N。
然后，可以通过扩展从 ct_ExpPack 恢复特殊密文。ExpPack的更多密码学细节在附录C中提供。

因此，ExpPack极大地减少了特殊密文的通信开销，代价是额外的扩展密钥（在我们的参数下高达5.16MiB）。然而，这些密钥可以存储在服务器上，并在同一客户端的多个查询中重用。

3.5.1 多用途扩展

ExpPack特别适合第3.4节中的协议，因为它可以用于我们统一模型中的多条路径（见图4）。对于双打包协议，扩展只执行一次，之后恢复的特殊密文可以重用于多个打包步骤。对于混合协议，我们利用 n < N 的事实，在 ct_ExpPack 中嵌入额外信息：我们同时嵌入用于打包的 s_i 值和用于后续多项式HE折叠步骤的索引信息。这使得单个扩展过程可以同时服务于打包（③）和多项式HE折叠（⑦）路径。

3.5.2 大规模并行性

ExpPack释放了Tiptoe风格打包中固有的大规模并行性，而这种并行性以前由于其高通信开销而无法实现。在扩展之后（通过并行化加速已在IVE [8]中得到广泛研究），接着是数千个独立的NTT和INT32 GEMM（见图2(b)），暴露了非常适合GPU的大规模并行性。

重新审视图2中的 B 批次SimplePIR，首先将 D₁×n 和 D₁×B 的SimplePIR输出通过沿着 D₁ 维度的每个长度-N 向量视为 ℛ_q 中的多项式（整数元素解释为系数）转换为多项式。这产生 ⌈D₁/N⌉×n 和 ⌈D₁/N⌉×B 个多项式，对其独立应用NTT。

设 H 和 G 表示NTT后的这两个多项式矩阵。接下来计算 H·C + ( 0 ∣ G )，其中 C 是一个对应于来自 B 个客户端的扩展特殊密文的 n×2B 多项式矩阵。由于RNS和NTT将多项式乘法转换为在 (#RNS·N) 个系数上的逐点乘法，这简化为 (#RNS·N)-并行的 ⌈D₁/N⌉×2B×n INT32 GEMM，并带有额外的模归约计算。

3.5.3 替代的环打包方法

与先前的环打包方法不同，ExpPack避免了限制并行性的顺序循环和导致额外内存开销的大量预计算。虽然我们的统一模型不排除在HintlessPIR [14]、YPIR [13]和InsPIRe [16]等研究中使用的替代打包方法，但许多方法在GPU上表现出固有的局限性。特别是，HintlessPIR和InsPIRe依赖于由数据依赖、有状态转换（例如，迭代旋转和累加）产生的长 O(n) 顺序循环，其中每一步依赖于前一步，限制了并行性。它们还依赖于大量的预计算，其结果占用大量内存；在我们对64GiB数据库的评估中，InsPIRe的CPU实现即使有256GiB的DRAM也耗尽了内存。YPIR避免了这种顺序结构，并表现出更接近ExpPack的计算模式，但没有提供第3.5.1节中相同的协议级优势。

3.6 将所有这些整合到我们的混合协议中

我们将通过针对组织为 N × 2 × … × 2 × D₀ = N × D_r × D₀ = 2¹² × 2⁸ × 2¹⁶ 结构（10亿条记录）的64GiB数据库的混合协议，逐步介绍我们的优化。

在离线阶段，服务器通过将数据库与一个 D₀×n 随机矩阵相乘来预计算提示，产生一个 (N×D_r)×n 的结果。然后它沿着 N 维度执行NTT以获得 H，一个 D_r×n 的多项式矩阵。每个多项式包含 3×N 个INT32元素（#RNS=3），占用48KiB。因此，H 占用15GiB的内存空间，成为协议中仅次于数据库的第二大数据对象。客户端在离线阶段也发送可重用的密钥（6.09MiB）。

在在线阶段，客户端发送一个长度-D₀ INT32向量（256KiB）用于标量HE折叠（①）和一个要扩展的RLWE密文（96KiB）。然后执行服务器端在线计算如下：

①（数据库GEMM）：查询向量与数据库（64GiB）之间的INT8-INT32 GEMM产生一个 N×D_r INT32矩阵。
③（ExpPack）：首先，扩展产生用于ExpPack的“特殊” C（n×2个多项式，120MiB）以及用于其余多项式HE折叠步骤的RGSW密文。将①的输出应用NTT，产生 D_r 个多项式（G，12MiB）。然后，执行INT32 GEMM以计算 H·C + ( 0 ∣ G ) 用于Tiptoe风格的打包，产生 D_r 个输出RLWE密文（24MiB）。
⑦（其余部分）：执行 log₂ D_r = 8 次RGSW-RLWE外积进行多项式HE折叠，以提取单个RLWE密文（96KiB）。在③之后，数据大小足够小，这些步骤对总成本的贡献很小。最终的模归约[37]产生一个32KiB的响应（一个具有两个多项式和#RNS=1的RLWE密文）。

4. VIPIR：一个GPU优化的PIR框架

我们基于统一模型开发了VIPIR，一个通用的GPU优化PIR框架。虽然它支持一系列PIR协议，但我们的目标不是覆盖所有现有变体，而是优化那些适合GPU执行的协议。通过将协议与GPU内核和执行策略协同设计，我们将PIR性能提升到实用水平，即使对于大型数据库也是如此。

PIR中的两个主要计算是GEMM和NTT。在我们的统一模型的主要执行路径中出现两种类型的GEMM：

• INT8-INT32 常规GEMM。这用于标量HE折叠（①）。由于我们像SimplePIR一样执行模 2³² 操作，模归约实际上是免费的。
• INT32 模GEMM。这用于多项式HE折叠（⑥）和环打包（③）。在这种情况下，我们对多项式进行操作，通过RNS分解和NTT实现 (#RNS·N)-并行GEMM。我们需要执行模 q_i 操作。

4.1 使用张量核心GEMM进行标量HE折叠

由于标量明文适合INT8，我们利用NVIDIA GPU张量核心[17]来利用其高计算吞吐量。例如，在我们评估中使用的H100 NVL GPU上，使用常规INT32操作的GEMM峰值吞吐量限制为30TOPS，而张量核心可实现高达1670TOPS的INT8操作吞吐量。

由于张量核心支持INT8 GEMM和INT32累加，我们将INT32矩阵分解为四个INT8矩阵。对于一个INT8矩阵 A 和一个INT32矩阵 B，我们将 B 分解为INT8矩阵 B₀, B₁, B₂, B₃，以计算 A·B，如下所示：
A·B₀ + 2⁸·(A·B₁) + 2¹⁶·(A·B₂) + 2²⁴·(A·B₃) mod 2³²，
其中每个 A·B_i 使用张量核心计算，乘以 {2⁸, 2¹⁶, 2²⁴} 通过移位实现，模 2³² 归约是免费的。

我们利用NVIDIA的prmt指令[40]（最初为视频处理引入）来高效执行矩阵分解。该指令从两个32位寄存器中选择四个任意字节，并将它们组装成一个32位目标寄存器。给定来自 B 的四个连续的INT32元素，我们使用两条prmt指令从每个元素中提取一个字节，并将它们打包成一个单一的32位寄存器，用于后续的张量核心操作。这使得分解可以使用仅 2×( B 的元素数) 条prmt指令实现。

所有这些操作，包括基于prmt的分解、INT8-INT8张量核心计算以及最终的基于移位的累加，都被融合到一个GPU内核中，以最小化内存开销。我们进一步利用最新NVIDIA GPU提供的高级GPU功能，例如异步全局到共享内存复制、向量化加载/存储指令、矩阵特定的共享内存加载（ldmatrix）和swizzling [41]。

我们的参数和协议选择旨在实现高效的张量核心执行。虽然到目前为止我们默认采用INT8明文和INT32密文，但这需要为标量HE设置明文模数为 2⁸（而不是SimplePIR的247-991范围）和密文模数为 2³²。相比之下，YPIR [13]和InsPIRe [16]中的打包方法不支持2的幂次密文模数，需要像多项式HE中那样进行模 q_i（q_i < 2³²）算术，这在张量核心执行下给累加和模归约带来了显著的开销[42]。

4.2 优化并行模GEMM

我们通过以下方式优化环打包和多项式HE折叠中使用的并行INT32模GEMM：（1）协调NTT和GEMM之间的布局不匹配，以及（2）减少模归约开销。

在ExpPack中，NTT [8]主导了扩展的计算，并倾向于使用 N-最内层布局，以实现每个多项式内的连续访问。然而，这种布局对于 (#RNS·N)-并行GEMM来说并不理想，在GEMM中，线程操作矩阵块，并且需要块内的连续访问以实现内存合并和共享内存重用[41]。为了弥补这种不匹配，我们在GEMM之前和之后插入布局转置，以在多边形主（NTT友好）和分块主（GEMM友好）布局之间进行转换。这实现了合并的全局内存访问、更高的共享内存重用和更大的GEMM块使用，从而提高了算术强度和吞吐量。

我们通过惰性归约[19, 18]进一步减少模归约开销。使用小于 2²⁹ 的RNS素数和64位累加器，我们将归约推迟最多32次乘加运算而不会溢出，从而分摊归约成本并降低指令开销。

4.3 数论变换（NTT）

我们通过采用单内核设计与Montgomery归约来优化NTT执行。对于 N=2¹² 的NTT，其工作集适合片上资源，我们可以用一个内核实现它，消除了内核间同步开销[43, 44]。我们还采用Montgomery归约[45]代替Barrett归约[46]，因为它更适合具有较少指令的GPU 32位整数数据路径[19]。

我们还融合内核以减少NTT与其相邻内核之间的内存流量。具体来说，在ExpPack中，我们将初始RNS分解与NTT融合（第3.5节），减少了临时结果的DRAM传输。我们还将模归约[37]与逆NTT融合。

4.4 内存容量感知的流式缓冲

尽管我们避免了数据库膨胀，但某些中间数据对象会占用大量内存空间。特别是，在我们的双打包协议中的第二次折叠将压缩的中间结果（每个客户端 8N×D₁ INT8）与标量HE密文（D₁×(n+1) INT32）相乘，产生每个客户端 8N×(n+1) 的结果（160MiB），对于64个批次总计10GiB；后续用于环打包的NTT进一步使此足迹增加三倍。

为了平衡批处理并行性和内存效率，我们引入了流式缓冲。对于 B 批次执行，我们将执行分区到CUDA流中，为每个流分配一个固定大小的缓冲区（大小为单个批次），并为每个流分配 B/#stream 个批次。每个流顺序处理其批次，重用缓冲区，将中间内存使用量从 O(B) 减少到 O(#stream)，同时保持并行性。我们选择性地将其应用于内存密集型阶段。

4.5 多GPU支持

为了进一步扩展吞吐量并支持更大的数据库大小，我们通过多GPU支持扩展了VIPIR。我们设计了两种并行化策略来分布数据和计算，两者都在单个协议中使用。

首先，由于数据库主导DRAM使用，我们将其分片到多个GPU上，以支持超过单个GPU内存容量的数据库大小。图5显示了SimplePIR的一个示例，其中每个GPU存储数据库的一个 (D₁/#GPU)×D₀ 分区，并与输入查询（一个 D₀×B 矩阵）独立执行GEMM，查询被广播到所有GPU。

图5：SimplePIR的双GPU在线计算流程。

在第一次折叠后，我们执行GPU间通信以过渡到批次分布，从而消除了后续阶段进一步通信的需求。如果第一次折叠产生一个 (D₁/#GPU)×B 的输出，我们执行一次全交换，使得每个GPU持有聚合结果的 D₁×(B/#GPU) 分区，有效地将每个GPU的批次大小减少到 B/#GPU。由于SimplePIR提示在所有查询和客户端之间共享，它必须驻留在每个GPU上以进行进一步的PIR步骤。因此，我们在提示预计算期间执行一次全收集。

表 1：不同数据库（DB）规模下无提示私有信息检索（PIR）实现的性能对比。​

前三个系统仅使用多项式同态加密（poly‑HE）折叠；接下来的六个系统基于 SimplePIR (S) 或 DoublePIR (D)（Henzinger 等，2023b）使用标量同态加密（scalar‑HE）折叠；最后两个系统则结合了 poly‑HE 与 scalar‑HE 折叠。绿色单元格表示每组中的最优值。我们报告的是单批次（B = 1）下的服务器端在线执行时间（毫秒）。每次查询的通信量（Comm./query）包含不可复用的密钥，但不包含可复用密钥（Reuse‑keys）。我们在 VIPIR 上使用了单张 H100 NVL 94GB GPU。OoM表示内存溢出（Out‑of‑Memory）失败。

5. 实现

VIPIR使用C++和CUDA实现。我们使用CUTLASS [47]中的CuTE布局抽象来实现INT8-INT32 GEMM，并利用NVIDIA集体通信库（NCCL）[48]进行高效的GPU间通信。我们在单批次执行中使用基于CUDA核心的单独内核进行INT8-INT32矩阵-向量乘法，因为这些操作是内存受限的，无法利用张量核心的吞吐量。VIPIR专注于加速服务器端PIR操作，这是主要的性能瓶颈；客户端操作使用基于CPU的HE库SEAL [49]实现，用于评估目的。

我们遵循标准实践[5, 15, 31, 25, 28]选择参数。对于标量HE，我们设置明文模数为 2⁸，密文模数为 2³²，在每次标量HE折叠步骤后归约到 2¹⁸。对于多项式HE（RLWE和RGSW），我们设置明文模数为 2¹⁸，密文模数为 q≃2⁸⁷（#RNS=3），可以通过模归约[37]归约到 q≃2²⁹（#RNS=1），其中为实现OnionPIRv2 [28]使用了稍大的参数。我们验证，当使用基于Lattice Estimator [50]的次数 n=1,280 和 N=2¹² 时，这些参数实现了128位安全性。参数配置在附录A和附录D中有更详细的说明。

6. 评估

6.1 实验设置

我们在多种PIR协议上测试了VIPIR，包括现有的——OnionPIRv2 [28]和SimplePIR/DoublePIR [31]；与先前工作[13, 15, 14]仅在环打包上不同的无提示协议——ExpPack(S)（基于SimplePIR：①→③）和ExpPack(D)（基于DoublePIR：①→②→①→③）；加上我们的双打包和混合协议。

我们主要在一个H100 NVL 94GB GPU上评估VIPIR的性能，其有效HBM带宽为3.48TB/s。为了展示多GPU可扩展性，我们还使用通过NVLink 3.0连接的两个H100 NVL 94GB GPU进行了实验，提供530.16GB/s的带宽。所有报告的执行时间仅反映服务器端在线执行时间，吞吐量以每秒查询数（QPS）衡量。当数据库包含 D 条1B记录时，我们为SimplePIR和DoublePIR（无论有无ExpPack）设置 D₀ ≈ D。对于双打包和混合协议，我们使用 D/D₀ = 2²⁰ 来最小化通信，除了64GiB数据库，我们默认使用 D/D₀ = 2¹⁹ 以避免大批次执行下的内存不足故障（表1报告单批次结果，因此仍使用 2²⁰）。

我们将VIPIR与最先进的PIR研究进行比较，包括Spiral [24]、OnionPIRv2 [28]、HintlessPIR [14]、YPIR [13]、InsPIRe [16]和SimplePIR/DoublePIR [31]。它们的开源CPU版本在具有Intel Xeon Gold 6348 CPU和256GiB DRAM的系统上执行。基线的配置在附录B中描述。

6.2 VIPIR与CPU基线对比

我们与先前CPU基线的广泛比较（表1）显示，VIPIR在所有基线中，在服务器端PIR计算方面实现了数个数量级的性能提升。与基线中表现最佳的基于DoublePIR的YPIR [13]相比，即使是我们最慢的配置ExpPack(D)也显示出12.5-77.9倍的改进。值得注意的是，性能增益随着数据库大小的增加而增加——对于1GiB为12.5倍，对于64GiB为77.9倍——突显了VIPIR对大型数据库的强大可扩展性。

由于我们内存容量感知的设计，对辅助数据对象的开销最小，VIPIR能够在有限的内存预算内高效地容纳大型数据库。相比之下，InsPIRe [16]和多项式HE折叠方法，如Spiral [24]和OnionPIRv2 [28]，即使在256GiB主机内存下也遇到内存不足（OoM）故障。然而，VIPIR尽管只有94GiB的可用内存，却能成功地在GPU上完全处理高达64GiB的数据库。

我们的新协议构建在保持有竞争力的服务器端计算性能的同时，展示了有希望的通信成本降低，即使有额外的计算开销。双打包和混合协议在其各自的协议组中实现了最低的每查询通信成本。同时，对于大型数据库，它们的服务器端性能与简单设计如ExpPack(S)相匹配。对于64GiB数据库，ExpPack(S)需要28.3毫秒的服务器计算，而双打包和混合协议分别需要33.0毫秒和33.5毫秒。

使用我们的协议，64GiB PIR的执行时间33.0-33.5毫秒，仅比数据库访问时间所施加的硬延迟界限高1.67-1.69倍。在我们评估的GPU系统上，从DRAM流式传输64GiB数据需要19.7毫秒。

VIPIR在执行现有协议方面也表现出色。与OnionPIRv2 [28]的CPU实现相比，VIPIR使用相同的协议实现了283-422倍的性能提升。虽然VIPIR目前的通信成本是CPU基线的两倍，但这可以通过用随机种子替换一半传输数据来平衡[38]，正如基线所做的那样。

表 2：基于 YPIR（Menon 和 Wu，2024）CPU 实现的普通 SimplePIR/DoublePIR（不支持批处理）与 32 批处理 VIPIR 的吞吐量（QPS）对比。

对多客户端批处理[8]的支持在评估吞吐量（QPS）时可以进一步扩大性能差距。虽然先前基于CPU的研究仅提供有限的批处理能力，但只要内存空间足够，VIPIR就完全支持批处理。如表2所示，在相同的SimplePIR/DoublePIR协议下，与单批次CPU实现[13]相比，VIPIR的32批次执行实现了188-1,029倍的更高QPS。

表 3：不同数据库规模下 VIPIR（H100 NVL 94GB，批处理大小 32）与先前基于 GPU 的 PIR 研究的吞吐量（QPS）对比。​

OoM表示内存溢出（Out‑of‑Memory）失败。

† 数值来自 Wang 等人（2025）的论文，基于 RTX 4090 显卡的报告结果。

6.3 VIPIR与GPU基线对比

也存在几项基于GPU的PIR研究。我们将VIPIR与PIRonGPU [7]（在同一GPU系统上执行）以及ShiftPIR [6]报告的性能进行了比较。PIRonGPU和ShiftPIR属于多项式HE折叠方法，由于NTT导致显著的数据库膨胀，限制了对大型数据库的可扩展性。ShiftPIR [6]试图通过分区数据库并将主机到设备传输与GPU计算重叠来缓解这一问题，但受到低主机-设备带宽的限制。

如表3所示，VIPIR在所有数据库大小和所有协议配置下，始终优于先前基于GPU的PIR实现。即使使用同样会导致数据库膨胀的OnionPIRv2，与PIRonGPU和ShiftPIR相比，VIPIR的QPS分别提高了65.1-75.0倍和41.3-49.6倍。此外，使用我们没有数据库膨胀的协议（表3中除OnionPIRv2外的所有协议）获得了更大的QPS增益，对于16GiB数据库达到175.7-405.1倍，对于64GiB达到652.6-931.3倍。

6.4 敏感性研究：执行参数

图6显示了吞吐量（QPS）如何随数据库大小、批次大小和协议变化。对于所有协议，QPS随着批次大小的增加而稳步增长。相对于单批次执行，32批次执行对ExpPack(S)、ExpPack(D)、双打包和混合协议的GPU吞吐量分别提高了3.69-6.62倍、11.33-15.06倍、2.64-4.95倍和3.33-6.02倍。ExpPack(D)特别强大的扩展源于其对环打包的最小依赖，环打包仅应用于 O(n²) 数据元素，因此与数据库维度无关。这意味着我们用于维度折叠的张量核心GEMM实现随批处理具有良好的扩展性。

随着数据库大小的增加，协议之间的QPS差距缩小，这使得我们的双打包和混合协议对大型数据库更具吸引力。以32批次执行为参考点，对于16GiB数据库，ExpPack(S)的QPS比双打包和混合协议分别高2.31倍和1.81倍。然而，当数据库大小增加到64GiB时，这些差距缩小到1.43倍和1.18倍。这是因为总执行时间越来越由第一次标量HE折叠步骤中执行的数据库GEMM主导，而扩展等步骤则是恒定开销。

图6还显示了每种情况下的峰值设备内存使用量。对于OnionPIR，即使是16GiB数据库的单批次执行，由于4倍的数据库膨胀，也已使用66.9GiB的内存。其他基于标量HE折叠的协议没有这个问题；即使对于64GiB数据库和高达32的批次大小，它们也只使用69.8-92.1GiB的内存。

图6：不同数据库大小、PIR协议和1到32批次大小下的每秒查询数（QPS）和峰值DRAM内存使用量。我们对VIPIR使用一个H100 NVL 94GB GPU。OoM指内存不足故障。

6.5 敏感性研究：GPU优化

图7显示了当我们逐步应用针对GEMM的GPU优化时，32批次执行时间的分解。对于64GiB数据库，如果没有张量核心（Base），初始数据库GEMM使用标准INT32操作，占服务器端在线计算总时间的60.5-74.0%。我们的张量核心数据库GEMM实现将其加速了3.89-4.05倍，将其份额减少到29.8-42.2%。应用转置INT32 GEMM进一步将第一次ExpPack计算中的GEMM加速了1.09-1.44倍。总之，这两种GEMM优化为16GiB数据库提供了1.44-1.53倍的端到端加速，为64GiB数据库提供了更大的2.04-2.29倍加速。

此外，由于我们高效的多GPU分区策略，VIPIR在使用两个GPU时实现了近乎完美的扩展，加速比为1.94-2.01倍，加速也来自每个GPU工作集的减少。虽然VIPIR支持更多GPU，但我们将它们的评估留作未来工作。

图7：随着我们的GPU优化——张量核心INT8-INT32 GEMM（+Tensor）、转置INT32 GEMM（+Transpose）和双GPU执行（+2-GPU）——逐步应用，在线服务器计算延迟。我们使用了通过NVLink连接的H100 NVL 94GB GPU。

6.6 混合协议的深入分析

确定数据库维度是协议构建的一个关键设计选择，对混合协议的影响尤其大。假设数据库总共包含 D 条记录。图8显示，改变 D₀ 在服务器计算和客户端-服务器通信之间创建了一个权衡。增加 D₀（减少 D/D₀）从初始折叠产生更压缩的输出（压缩因子：4(n+1)/D₀ 倍），减少了环打包的数据量。相比之下，D₀ 对数据库GEMM时间的影响可以忽略不计，因为对于批次大小 B，无论 D₀ 如何，(D/D₀)×B×D₀ GEMM 都涉及相同数量的 D·B 次乘加运算。因此，增加 D₀ 会减少服务器总执行时间，但会增加通信量（O(D₀)）以发送沿 D₀ 维度的索引的加密独热表示。

图8：在H100 NVL 94GB GPU上，混合协议中延迟分解和在线通信量随 D/D₀ 的变化。OoM指内存不足故障。

我们在实验中主要使用 D/D₀ = 2²⁰ 来最小化服务器-客户端通信，因为我们的优化在很大程度上减轻了计算成本。如果我们选择通信量更大的设置，如 D/D₀ = 2¹⁷，对于小于64GiB的数据库，我们可以获得1.59-1.76倍的加速（对比 2²⁰），对于64GiB获得1.21倍的加速（对比 2¹⁹）。我们利用这种权衡在图1中展示了多个数据点。

7. 相关工作

自从Chor等人的开创性工作[1]以来，已经提出了许多PIR协议。XPIR [23]采用单次多项式HE折叠步骤。SealPIR [21]引入了查询扩展，并通过重新解释步骤[51, 52]执行多次多项式HE折叠。WhisPIR [27]随后引入了一种扩展方法，降低了扩展密钥的开销。在SealPIR的基础上，MulPIR [38]使用RLWE密文间乘法，而OnionPIR [25]通过使用RGSW-RLWE外积[39]进行非初始折叠步骤，进一步提高了效率。OnionPIRv2 [28]通过硬件友好的实现和协议微调进一步优化了设计。Spiral [24]探索用MLWE [35]代替RLWE，Respire [22]利用基于子环的RLWE降维来处理小记录数据库。KsPIR [20]则提出使用加密的RLWE矩阵-向量乘法进行多项式HE折叠。

与此同时，自SimplePIR [31]引入以来，标量HE维度折叠已获得关注。SimplePIR通过要求客户端下载一个大的可重用提示来减少通信。然而，这个提示必须在每次数据库更新时更换，并且需要大量的客户端存储。无提示协议，包括Tiptoe [15]、HintlessPIR [14]和YPIR [13]，对SimplePIR应用各种环打包方法，这些方法建立在其他背景下使用的打包技术[12, 9, 11]之上。可能存在更多适用于PIR的打包方法，例如HERMES [10]。InsPIRe [16]提出了一种新的打包方法，利用大量预计算来降低在线计算成本，并引入了一种将数据库条目插值到函数中的多项式HE折叠技术。

尽管在改进PIR协议方面做了大量工作，但专注于其硬件加速的研究相对较少。PIRonGPU [7]使用HEonGPU库[53]提供了SealPIR [21]的开源GPU实现。ShiftPIR [6]提出了一种GPU感知的协议，通过将部分查询生成卸载到服务器来减少客户端开销。Lam等人[54]和CIP-PIR [55]探索了多服务器PIR协议的GPU加速，这些协议依赖于我们认为是限制性的非共谋假设，因此在本工作中不予涉及。

最近也有一些针对PIR加速的架构提案。IVE [8]提出了一种ASIC加速器设计，专门用于处理OnionPIR查询批次。INSPIRE [56]（不同于InsPIRe）为SSD上的PIR引入了一种存内加速器。SmartPIR [57]通过将HE操作卸载到计算存储设备（CSD）扩展了这种方法。Conflux [58]类似地利用CSD进行关键词PIR，这是基于索引的PIR的扩展。

8. 结论

我们推出了VIPIR，一个通用的GPU框架，通过统一的、以GPU为中心的视角重新思考PIR的设计和实现。通过引入统一的分析模型，我们揭示了先前PIR协议中的基本权衡，特别是密文膨胀与数据库（明文）膨胀之间的矛盾。在这些见解的指导下，我们提出了双打包和混合协议，它们在不过度要求内存容量的情况下提高了通信效率。一个关键因素是ExpPack，一种新颖的环打包技术，它在降低通信成本的同时，暴露了非常适合GPU的大规模并行性。VIPIR还包括各种架构感知的优化，包括张量核心加速的GEMM、优化的NTT执行、内存高效的调度和多GPU支持。广泛的评估表明，与先前基于CPU和GPU的PIR系统相比，VIPIR实现了数个数量级的性能提升，同时降低了通信成本并避免了内存不足故障。