NVIDIA CUDA 13.3 Adds Hardware Carryless Multiplication Boosting Cryptography Throughput
NVIDIA CUDA 13.3 introduces the clmad PTX instruction for hardware-accelerated carryless multiplication on Ampere and newer GPUs. It delivers up to 18.8x speedups for GHASH (AES-GCM) on Blackwell B200 and accelerates zero-knowledge proof protocols.
Impact: High
Why it matters
Carryless multiplication is a fundamental building block for modern cryptography, error-correction codes, and zero-knowledge proofs. Native GPU support eliminates slow bitsliced emulation, dramatically increasing throughput for secure communication and ZK-proving pipelines.
TL;DR
- 01NVIDIA CUDA 13.3 introduces the clmad PTX instruction, enabling hardware-accelerated carryless multiplication on Ampere and newer GPUs.
- 02GHASH throughput for AES-GCM jumps up to 18.8x on the NVIDIA B200, reaching 6.3 TB/s.
- 03Zero-knowledge proof systems benefit from a 3x to 13x acceleration in sum-check protocols over GF(2^128).
Key facts
- Minimum GPU Support
- Ampere (SM 80) and newer
- Peak GHASH B200 Throughput
- 6.3 TB/s
- Peak GHASH RTX 5090 Throughput
- 1,300 GB/s
Hardware-Accelerated Carryless Multiplication
The new clmad PTX instruction computes the carryless multiplication of two 64-bit inputs into a 128-bit result. It is exposed in CUDA 13.3 via .hi and .lo variants to calculate the respective halves of the result with the addition of a 64-bit accumulator.
Example PTX Implementation
Developers can write inline PTX assembly straight in their GPU kernels for quick binary field operations:
__device__ inline uint128_t clmad_mul_128(uint64_t a, uint64_t b, uint128_t acc) {
uint64_t acc_hi = (uint64_t)(acc >> 64);
uint64_t acc_lo = (uint64_t)acc;
uint64_t lo, hi;
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;
}Benchmarks and Real-World Impact
Performance metrics for key cryptographic operations show massive gains over traditional bitsliced implementations:
- GHASH on B200: Up to
18.8xspeedup, achieving6.3 TB/sthroughput. - GHASH on RTX 5090: Reaches
1,300 GB/sthroughput, representing a2ximprovement. - Sum-Check over GF(2^128): Commonly used in zero-knowledge (ZK) proving systems, achieving a
3xto13xacceleration usingclmad.
Try it in 2 minutes
__device__ inline uint128_t clmad_mul_128(uint64_t a, uint64_t b, uint128_t acc) {
uint64_t acc_hi = (uint64_t)(acc >> 64);
uint64_t acc_lo = (uint64_t)acc;
uint64_t lo, hi;
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;
}cpp
✓ When to use
- When optimizing GHASH, AES-GCM, or zero-knowledge proof protocols on Ampere or newer GPUs.
- When implementing Reed-Solomon, BCH, or quantum stabilizer codes on modern NVIDIA hardware.
✕ When NOT to use
- When targeting GPUs older than the Ampere architecture (pre-SM 80).
- When the cryptographic or error-correction algorithm does not rely on binary field arithmetic.
What to do today
- Download CUDA 13.3 to access the new clmad PTX instruction.
- Integrate clmad into binary field multiplication kernels for cryptosystems like AES-GCM or zero-knowledge proofs.
Sources