网络安全

在 NVIDIA CUDA 13.3 中使用无带乘法构建更快的加密技术

2026年 7月 15日

作者：Jean-Luc Watson, Charles Gouert, Michael B. Sullivan 和 G. Edward Suh

15 年多来，x86 CPU 一直附带用于无携带乘法的专用硬件指令。它是一个小但却很棘手的基元，位于经过身份验证的加密、纠错码和现代零知识证明的背后。

到目前为止，NVIDIA GPU 缺乏对该操作的原生支持。NVIDIA CUDA 13.3 通过 clmad 弥补了这一差距。clmad3 是所有 NVIDIA Ampere 及更新 GPU (SM 80+) 上可用的新 PTX 指令。在本文中，我们展示了当无携带乘法最终成为硬件加速的 GPU 基元时的可能性。

我们对两种依赖它的加密工作负载进行了基准测试：GHASH、AES-GCM ( TLS、VPN 和大多数数据中心加密背后的 AEAD 密码) 内的完整性哈希，以及总和检查协议 (先进零知识验证系统的主要内环) 。在 NVIDIA B200 上，GHASH 吞吐量达到约 6.3 TB/s，接近 DRAM 读取带宽，比之前的比特片技术快 18.8 倍。对 \(GF(2^{128})\) 的求和检查可将先前的技术水平提高 4 – 13%。

为什么这不仅仅是 AES-GCM 的重要意义？在种类繁多的加密和编码理论工作负载下，无携带乘法是共享内核。用于存储系统和电信基带处理的 CRC 和 Reed–Solomon 代码、用于闪存的 BCH 代码、量子稳定器代码、几种后量子加密方案，以及支撑现代零知识证明系统 (例如 Binius) 的二进制场算法。GPU 上的硬件加速改变了所有这些工作负载 (在每个已部署的 Ampere 或更高版本系统上) 的成本结构。

本文的目标受众是将加密技术集成到 GPU 工作流中的 CUDA 开发者，以及基于二进制域协议的安全和隐私研究人员。继续下载 CUDA 13.3；下面的内联 PTX 示例可以直接进入您的内核。

二进制扩展程序字段

最小的有限域 \(GF(2)\) 只是一个位，其中加法是 XOR，乘法是 AND。在加密用例中，位通常组合成 二进制扩展字段 \(GF(2^n)\)，其中 \(n\) 位表示多项式在 \(GF(2)\) 上的二进制系数，而运算则以模的形式取不可约多项式。

该领域的多项式加法仍然是简单的 XOR (元素级系数加法) ，但乘法需要在每个输入的位之间进行长乘法运算。

具体来说，输入 \(a\) 和 \(b\) 的扩展场乘法的第 2 个所得系数是 \(\bigoplus_{j=0}^{i} a_j b_{i-j}\) — — 单个位积的 XOR，然后是该字段的不可约多项式的缩减。

例如，2 位扩展字段 \(GF(2^2)\) 可能有一个不可约多项式，即 \(X^2 + X + 1\)。

可以将此计算模拟为一系列 AND 和 XOR 运算，但提取单个位值所需的重复位移和掩码会导致大量开销。x86 提供了对这类场乘法的直接硬件支持，主要支持 AES-GCM。然而，到目前为止，NVIDIA GPU 尚未实现无携带乘法，因此开发者需要采用位切片电路等替代技术来提高性能。

硬件支持无携带乘法

CUDA 13.3 为无携带乘积累加指令 clmad 添加了新的 PTX 支持。

与 x86 中的 PCLMULQDQ 非常相似，clmad 可将两个 64 位输入乘以 128 位结果。.hi 和 .lo 变体通过添加 64 位累加器，分别计算 128 位输出的上半部分和下半部分，如下所示：

__device__ inline uint128_t clmad_mul_128(uint64_t a, uint64_t b, uint128_t acc) {
	uint64_t acc_lo = (uint64_t)acc;
	uint64_t acc_hi = (uint64_t)(acc >> 64);

	uint64_t lo, hi;
	
	// Inline PTX: calculate lower and higher 64 bits of result
	asm("clmad.lo.u64 %0, %1, %2, %3;" : "=l"(lo) : "l"(a), "l"(b), "l"(acc_lo));
	asm("clmad.hi.u64 %0, %1, %2, %3;" : "=l"(hi) : "l"(a), "l"(b), "l"(acc_hi));
	
	return ((uint128_t)hi << 64) | lo;
}

从 Ampere ( sm_80 或更高版本) 开始，此指令在 NVIDIA GPU 上的硬件中进行加速，以便在现代 NVIDIA GPU 上实现快速的二进制扩展域乘法。

在本文的其余部分，我们来看看 clmad 如何加速两个加密用例 — — GHASH 和 sumcheck 协议，这两个用例都在更大的 \(GF(2^{128})\) 字段中运行。在该字段中，可以使用 Karatsuba 算法在 6 个 clmad 指令中对元素进行乘法运算：

__device__ inline uint256_t clmad_mul_256(uint128_t a, uint128_t b) {
	// Isolate low and high halves of inputs
	uint64_t a0 = (a & 0xFFFFFFFFFFFFFFFFull), a1 = (a >> 64);
	uint64_t b0 = (b & 0xFFFFFFFFFFFFFFFFull), b1 = (b >> 64);

	// Karatsuba: divide-and-conquer with smaller 64x64->128 bit carryless mults
	uint128_t z0 = clmad_mul_128(a0, b0, 0);
	uint128_t z1 = clmad_mul_128(a0 ^ a1, b0 ^ b1, 0);
	uint128_t z2 = clmad_mul_128(a1, b1, 0);
	z1 ^= z0 ^ z2;
	z0 ^= z1 << 64;
	z2 ^= z1 >> 64;
	return {z2, z0}; // to be reduced back to GF(2^128)
}

要从 256 位结果中返回到 128 位字段，可以通过基于乘法的 Barrett 归约与其他 CLMAD 调用来计算多项式归约，也可以通过模拟多项式长除法的直接 Shift-and-XOR 循环来计算多项式归约。

GHASH 加速

GHASH 非常适合 clmad，因为它在很大程度上基于 \(GF(2^{128})\) 中的二进制乘法，模拟了不可约多项式 \(X^{128} + X^7 + X^2 + X + 1\)。

在更广泛的 AES-GCM 背景下，GHASH 计算所有输入数据 (即密码、其他经过身份验证的数据和最终长度块) 的核心身份验证哈希。GHASH 将输入拆分成 128 位块，对于每个块，将其 XOR 成一个正在运行的累加器，并在 \(GF(2^{128})\) 中将累加器乘以一个哈希键 (H) (由使用 AES 加密零块衍生而来) 。

这种乘法运算可以通过 clmad 和模块化归约来实现。处理完所有块后，结果是一个 128 位的身份验证哈希。然后，使用初始计数器块的 AES 加密对该哈希进行 XOR 运算，从而形成身份验证标记。

在 NVIDIA GeForce RTX 5090 D 上，借助 clmad，GHASH 的峰值吞吐量约为 1300 GB/s，与通过比特切片实现的 GHASH 相比，峰值吞吐量提高了 2 倍。

相反，在 B200 上，我们观察到峰值吞吐量约为 6335 GB/s (接近 DRAM 读取带宽) ，吞吐量比位切片版本高 18.8 倍。在比特片基准测试中，B200 实际上比 RTX 5090 慢 ( SM 较少，时钟频率较低) ，但 CLMAD 的硬件加速使其 GHASH 速度非常快。

零知识：求和检查加速

二进制扩展域用于零知识 (ZK) 验证协议。它们使一方能够执行计算，然后以加密方式向另一方证明结果是正确的，而无需验证器重做工作。二进制字段在证明算法的结果时很有帮助，这些算法会自然地作为逐位运算进行评估，例如哈希和模拟常规整数运算。当然，ZK 协议需要进行多次无携带乘法运算，如证明协议中所述。

sum-check 协议是构成许多验证系统基础的一个基本基元。它使一方能够证明对布尔输入 (即 \(\Sigma_{b_1,b_2,\ldots,b_n \in \{0, 1\}}g(b_1, b_2, …, b_n)\)) 的 \(n\) 变量多项式求和的了解程度，以便验证器的工作在 \(n\) 中是线性的，并且只需要在一个随机点评估 \(g\)。

Sum-check 在 \(n\) 轮中执行，每个变量一轮。我们专注于处理多项式合成 (\(g(\cdot) = c_1(\cdot) \times c_2(\cdot) \times \ldots \times c_d(\cdot)\)) 的简单实例化，其中每个 \(c_i\) 由 \(GF(2^{128})\) 中的 \(2^n\) 评估列表表示。在 \(r\) 回合中，验证者：

计算已声明的评估的乘积和 (在每个 \(2^r\) 评估行的多项式上合成乘积) ，
插补 \(d\) 点的多项式，并计算其积和，以生成该回合的一元多项式，最后
在验证器选择的随机点插入多项式，并通过 \(2^{r-1}\) 评估从一组减半的多项式开始下一轮。

每个插值和乘积都依赖于多项式评估之间的大量二进制扩展域乘法 (根据输入和合成大小进行缩放) ，并且几乎完全可并行！

我们将 \(GF(2^{128})\) 中基于 clmad 的线性时间线性空间和检查协议与 Irreducible 为相同的现场操作开发的位切片电路进行比较。在 RTX 5090 D 上使用硬件支持二进制字段，可将求和检查性能提升 3-4 倍，在 B200 GPU 上可提升高达 13 倍，并且随着多项式和合成大小的增加，性能优势会略有提升。其中，Irreducible对应的中文部分为”Irreducible”，保持原有链接和标签。

总结

CUDA 13.3 为 NVIDIA GPU 带来了一个缺失的基元：硬件加速的无携带乘法。我们对两个工作负载 ( GHASH 和 sum-check 协议) 进行了基准测试，在 B200 上，速度分别提高了 18.8 倍和 13 倍。最大的收益显示了现代密码学的发展方向：大输入身份验证加密和二进制域零知识证明。

开始使用：

下载 CUDA 13.3
在您自己的内核中试用上述内联 PTX 示例 (需要 Ampere 或更高版本)
针对加密工作负载？请参阅 cuPQC SDK。
构建电信应用？请参阅 NVIDIA Aerial 文档。

关于作者

Jean-Luc Watson 是 NVIDIA 安全和隐私研究组的研究科学家。他的研究旨在通过应用加密和系统设计的交叉工作，使加速应用程序更加隐私和安全。此前，他在加州大学伯克利分校攻读博士学位，专注于基于 GPU 的安全计算以及适用于嵌入式系统的隐私优先协议。

查看 Jean-Luc Watson 所有文章

Charles Gouert 是 NVIDIA 编程系统和应用研究团队的研究科学家。他的专业领域是应用加密、安全且可验证的计算，以及加密基元的硬件加速，例如全同态加密，这是他在特拉华大学博士论文的重点。他拥有理科学士学位。来自特拉华大学计算机工程系。

查看 Charles Gouert 所有文章

Michael B. Sullivan 是 NVIDIA 的首席研究科学家。他获得了学士学位。计算机工程专业，学士学位。数学专业，以及理科硕士学位。计算机科学专业，来自乔治・梅森大学和 M.S.E.。德克萨斯大学奥斯汀分校计算机架构专业博士。他的研究重点是一个高效、可靠且安全的计算机系统，包括低成本的安全性、内存安全、强大的内存系统可靠性，以及用于特定应用加速的软硬件技术。

查看 Michael B. Sullivan 所有文章

G. Edward Suh 是安全和隐私研究小组的高级研究总监。2007 年至 2023 年，他还是康奈尔大学电气与计算机工程学院的客座教授。他获得了学士学位。首尔大学电气工程专业，理科硕士学位。以及麻省理工学院 (MIT) 电气工程和计算机科学博士学位。他最近的研究重点是构建安全的计算系统，以实现安全和隐私的 AI，并使用 AI 来提高计算机系统的安全性。他是 ACM 和 IEEE 的研究员。

查看 G. Edward Suh 所有文章