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 倍。对 𝐺⁡𝐹⁡(2128) 的求和检查可将先前的技术水平提高 4 – 13%。

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

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

二进制扩展程序字段

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

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

具体来说，输入 𝑎 和 𝑏 的扩展场乘法的第 2 个所得系数是 ⨁𝑖𝑗=0𝑎𝑗⁢𝑏𝑖−𝑗 — — 单个位积的 XOR，然后是该字段的不可约多项式的缩减。

例如，2 位扩展字段 𝐺⁡𝐹⁡(22) 可能有一个不可约多项式，即 𝑋2 +𝑋 +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 协议，这两个用例都在更大的 𝐺⁡𝐹⁡(2128) 字段中运行。在该字段中，可以使用 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，因为它在很大程度上基于 𝐺⁡𝐹⁡(2128) 中的二进制乘法，模拟了不可约多项式 𝑋128 +𝑋7 +𝑋2 +𝑋 +1。

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

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

图 1. GHASH 吞吐量测量值，单位为  𝐺⁡𝐹⁡(2128) ，用于长度不超过 2 GB 的输入有效载荷，并将使用 ALU 运算的位切片电路与 B200 和 NVIDIA GeForce RTX 5090 D 上基于 CLMAD 的版本进行比较

在 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 协议是构成许多验证系统基础的一个基本基元。它使一方能够证明对布尔输入 (即 Σ𝑏1,𝑏2,…,𝑏𝑛∈{0,1}⁢𝑔⁡(𝑏1,𝑏2,…,𝑏𝑛)) 的 𝑛 变量多项式求和的了解程度，以便验证器的工作在 𝑛 中是线性的，并且只需要在一个随机点评估 𝑔。

Sum-check 在 𝑛 轮中执行，每个变量一轮。我们专注于处理多项式合成 (𝑔⁡( ⋅) =𝑐1⁡( ⋅) ×𝑐2⁡( ⋅) ×… ×𝑐𝑑⁡( ⋅)) 的简单实例化，其中每个 𝑐𝑖 由 𝐺⁡𝐹⁡(2128) 中的 2𝑛 评估列表表示。在 𝑟 回合中，验证者：

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

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

图 2. 在各种多项式大小和合成大小下，对 𝐺⁡𝐹⁡(2128)  中的协议延迟进行求和检查，将使用 ALU 运算的位切片电路与基于 CLMAD 的版本进行比较

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

总结

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