Building Faster Cryptography with Carryless Multiplication in NVIDIA CUDA 13.3

For over fifteen years, x86 CPUs have shipped with a dedicated hardware instruction for carryless multiplication. It’s a small but stubborn primitive that sits underneath authenticated encryption, error-correcting codes, and modern zero-knowledge proofs. 

Until now, NVIDIA GPUs lacked native support for this operation. NVIDIA CUDA 13.3 closes that gap with clmad, a new PTX instruction available on all NVIDIA Ampere and newer GPUs (SM 80+). In this post, we show what is possible when carryless multiplication is finally a hardware-accelerated GPU primitive. 

We benchmark two cryptographic workloads that depend on it: GHASH, the integrity hash inside AES-GCM (the AEAD cipher behind TLS, VPN, and most data-center encryption), and the sum-check protocol, the workhorse inner loop of advanced zero-knowledge proving systems. On the NVIDIA B200, GHASH throughput reaches ~6.3 TB/s—close to DRAM read bandwidth and up to 18.8x faster than the prior bitsliced state of the art. Sum-check over \(GF(2^{128})\) speeds up the prior state of the art by 4–13×.

Why does this matter beyond AES-GCM? Carryless multiplication is the shared kernel underneath a surprisingly wide range of cryptographic and coding-theoretic workloads. CRC and Reed–Solomon codes used in storage systems and telecom baseband processing, BCH codes for flash memory, quantum stabilizer codes, several post-quantum cryptographic schemes, and the binary-field arithmetic that underpins modern zero-knowledge proving systems such as Binius. Hardware acceleration on the GPU changes the cost structure for all these workloads, on every Ampere-or-later system already deployed. 

This post is written for CUDA developers integrating cryptography into GPU pipelines and for security and privacy researchers building on binary-field protocols. Download CUDA 13.3 to follow along; the inline PTX examples below can drop straight into your kernels.

Binary extension fields

The smallest finite field, \(GF(2)\), is just a single bit, where additions are XORs and multiplications are ANDs. In cryptographic use cases, bits are typically combined into binary extension fields, \(GF(2^n)\), where \(n\) bits represent the binary coefficients of a polynomial over \(GF(2)\) and operations are taken modulo an irreducible polynomial. 

Polynomial addition in the field is still a simple XOR (element-wise coefficient addition), but multiplication requires a long multiplication between the bits of each input. 

Specifically, the \(i\)’th resulting coefficient of an extension-field multiplication of inputs \(a\) and \(b\) is \(\bigoplus_{j=0}^{i} a_j b_{i-j}\) – the XOR of individual bit products, followed by a reduction by the field’s irreducible polynomial. 

For example, the 2-bit extension field \(GF(2^2)\) has one possible irreducible polynomial, \(X^2 + X + 1\).

It’s possible to emulate this computation as a series of AND and XOR operations, but the repeated bitshifts and masking required to extract individual bit values result in significant overhead. Direct hardware support for this type of field multiplication has been available in x86, primarily supporting AES-GCM. However, until now, carryless multiplication hasn’t been available on NVIDIA GPUs, requiring developers to adopt alternative techniques like bitsliced circuits to improve performance.

Hardware support for carryless multiplication

CUDA 13.3 adds new PTX support for a carryless multiply-accumulate instruction, clmad.

Much like PCLMULQDQ in x86, clmad performs a carryless multiplication of two 64-bit inputs into a 128-bit result. .hi and .lo variants calculate the top and bottom halves, respectively, of the 128-bit output with the addition of a 64-bit accumulator, as shown:

__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;
}

This instruction is accelerated in hardware on NVIDIA GPUs since Ampere (sm_80 or higher) for quick binary extension field multiplication on modern NVIDIA GPUs.

In the remainder of this post, we look at how clmad can accelerate two cryptographic use cases – GHASH and the sumcheck protocol–both of which operate in the larger \(GF(2^{128})\) field. In that field, elements can be multiplied using the Karatsuba algorithm in 6 clmad instructions:

__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)
}

To get back to a 128-bit field from the 256-bit result, polynomial reduction can be computed either by a multiplication-based Barrett reduction with additional CLMAD calls or by a direct shift-and-XOR loop that emulates polynomial long-division.

GHASH acceleration

GHASH is a great candidate for clmad because it’s heavily based on binary multiplication in \(GF(2^{128})\), modulo the irreducible polynomial \(X^{128} + X^7 + X^2 + X + 1\).

In the broader context of AES-GCM, GHASH computes the core authentication hash over all of the input data (i.e., ciphertext, additional authenticated data, and a final length block). GHASH splits the input into 128-bit blocks and, for each block, XORs it into a running accumulator and multiplies the accumulator by a hash key (H) — derived from encrypting a zero block with AES — in \(GF(2^{128})\). 

This multiplication operation can be achieved with clmad plus a modular reduction. After all blocks are processed, the result is a 128-bit authentication hash. The authentication tag is then formed by XORing this hash with an AES encryption of the initial counter block.

On an NVIDIA GeForce RTX 5090, we achieve a peak throughput of ~1,300 GB/s for GHASH with clmad, which results in a 2x higher peak throughput relative to GHASH implemented with bitslicing. 

Conversely, on a B200, we observe a peak throughput of ~6,335 GB/s (close to the DRAM read bandwidth), which exhibits up to 18.8x higher throughput than the bitsliced variant. On the bitsliced baseline, the B200 is actually slower than the RTX 5090 (fewer SMs, lower clock), but CLMAD’s hardware acceleration makes its GHASH very fast.

Zero-knowledge: Sum-check acceleration

Binary extension fields are being used in zero-knowledge (ZK) proving protocols. They enable one party to perform a computation and then cryptographically prove to another party that the result is correct—without requiring the verifier to redo the work. Binary fields are helpful when proving results for algorithms that are naturally evaluated as bitwise operations, like hashing and emulating regular integer operations. Naturally, ZK protocols require many carryless multiplications.

A fundamental primitive that forms the basis for many proving systems is the sum-check protocol. It enables a party to prove knowledge of a \(n\)-variate polynomial’s sum over boolean inputs (that is, \(\Sigma_{b_1,b_2,\ldots,b_n \in \{0, 1\}}g(b_1, b_2, …, b_n)\)) so that the verifier’s work is just linear in \(n\) and only needs to evaluate \(g\) at one random point.

Sum-check operates in \(n\) rounds, one for each variable. We focus on a simple instantiation that processes a composition of polynomials (\(g(\cdot) = c_1(\cdot) \times c_2(\cdot) \times \ldots \times c_d(\cdot)\)), where each \(c_i\) is represented by a list of \(2^n\) evaluations in \(GF(2^{128})\). At round \(r\), the prover:

1. Calculates the claimed product sum of the evaluations (composing products across polynomials at each of \(2^r\) evaluation rows),
2. Interpolates the polynomials at \(d\) points and computes their product sum to yield a univariate polynomial for the round, and finally
3. Interpolates the polynomials at a random point chosen by the verifier, starting the next round with a halved set of polynomials with \(2^{r-1}\) evaluations.

Each interpolation and product relies on a significant number of binary extension field multiplications between polynomial evaluations – scaling with the input and composition sizes – and is almost entirely parallelizable!

We compare a linear-time, linear-space sum-check protocol with clmad-based field operations for \(GF(2^{128})\) to one using the bitsliced circuit developed by Irreducible for the same field operations. Using hardware support for binary fields improves sum-check performance by 3-4x on an RTX 5090 and up to 13x on a B200 GPU, with performance benefit growing slightly as the polynomial and composition sizes increase.

Conclusion

CUDA 13.3 brings a missing primitive to NVIDIA GPUs: hardware-accelerated carryless multiplication. The two workloads we benchmarked—GHASH and the sum-check protocol—see speedups of up to 18.8x and 13x, respectively, on the B200. The largest gains show up exactly where modern cryptography is heading: large-input authenticated encryption and binary-field zero-knowledge proofs.

Get started:

• Download CUDA 13.3 
• Try the inline PTX examples above in your own kernels (Ampere or later required)
• Targeting cryptographic workloads? See the cuPQC SDK.
• Building telco applications? See the NVIDIA Aerial documentation.