Skip to content

GPU-accelerated
cryptographic
math libraries

High-performance, low-level GPU math libraries for developers building both classical and next-generation cryptographic applications.

Get Started

Overview

cuPQC is a CUDA-based SDK that accelerates cryptographic primitives on the GPU. It gives you low-level, high-performance building blocks you assemble into the protocols and schemes your application needs.

Use it as a fast backend inside an existing library, or as the foundation for new GPU-native crypto code across NVIDIA GPUs.

Highlights
  • Device extension: call math primitives directly from CUDA kernels and fuse them into a single device kernel, without host round-trips
  • Unified, modular API: reusable components that compose across domains and parameter sets
  • Link-time optimization: selects domain- and parameter-tuned GPU kernels at link time
  • Broad GPU coverage: runs consistently from edge GPUs (Jetson) to data center GPUs
SDK Architecture
Apps End-user applications and services
Crypto schemes ClassicalPQCZKPFHE & more
cuPQC SDK cuPQC-BigInt cuPQC-NTT cuPQC-Hash cuPQC-PK
Runtime CUDA
Hardware NVIDIA GPU

Build cryptography on the GPU

Design and implement cryptographic schemes by composing GPU primitives inside your CUDA kernels. One composable operator model, from low-level math to the protocol logic you're building.

#include <bigint.hpp>
using namespace cupqc;

// One 256-bit integer per thread.
using BI256 = decltype(BitWidth<256>() + SM<800>() + Thread());

__global__ void mul_mod_kernel(uint32_t* products,
                               const uint32_t* buf_a,
                               const uint32_t* buf_b,
                               const uint32_t* buf_m) {
    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;

    const BI256::bigint a(buf_a, i);   // load instance i
    const BI256::bigint b(buf_b, i);
    const BI256::bigint m(buf_m, i);

    // (a * b) mod m, entirely in your own kernel.
    a.mul_mod(b, m).store(products, i);
}

#include <ntt.hpp>
using namespace cupqc;

// 1024-point forward transform, one polynomial per block.
using NTT = decltype(Algorithm<algorithm::NTT>()
          + Direction<nttDirection::FORWARD>()
          + Precision<uint32_t>() + Size<1024>()
          + Block() + BlockDim<128>());

__global__ void ntt_kernel(uint32_t* polys,
                           const uint32_t* twiddles,
                           const nttConst<uint32_t> ntt_const) {
    uint32_t* poly = polys + blockIdx.x * NTT::Size;
    extern __shared__ uint32_t sdata[];

    NTT().load_to_mont(sdata, poly, ntt_const);
    __syncthreads();
    NTT().execute(sdata, twiddles, ntt_const.p);
    __syncthreads();
    NTT().store_from_mont(sdata, poly, ntt_const);
}

#include <hash.hpp>
using namespace cupqc;

using SHA256 = decltype(SHA2_256() + Thread());

__global__ void hash_kernel(uint8_t* digest,
                            const uint8_t* msg, size_t len) {
    SHA256 hash {};
    if (threadIdx.x == 0) {
        hash.reset();
        hash.update(msg, len);
        hash.finalize();
        hash.digest(digest, SHA256::digest_size);
    }
}

Inside cuPQC SDK

Big Integer
cuPQC-BigInt

Fixed-width multi-precision integer arithmetic

32–4096 bitsModularBarrettMontgomery
Transforms
cuPQC-NTT

Number Theoretic Transform for polynomial arithmetic

ForwardInverseBabyBearKoalaBearCustom primes
Hashing
cuPQC-Hash

Hash functions and Merkle trees

SHA-2SHA-3SHAKEPoseidon2
Public Key
cuPQC-PK

Public-Key Cryptography

ML-KEMML-DSA

Release notes

View all release notes →
0.6.0 New cuPQC-BigInt and cuPQC-NTT libraries for multi-precision integer arithmetic and number theoretic transforms
0.4.1 Library renaming, Poseidon2 KoalaBear field, and multi-block Merkle trees
0.4.0 Poseidon2 hash function and the Merkle Tree API

These are brief summaries. See the full release notes for complete details.