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-PK cuPQC-Hash
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 <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);
    }
}

#include <pk.hpp>
using namespace cupqc;

using MLKEM512Key = decltype(ML_KEM_512()
                  + Function<function::Keygen>()
                  + Block() + BlockDim<128>());

__global__ void keygen_kernel(uint8_t* public_keys,
                              uint8_t* secret_keys,
                              uint8_t* workspace,
                              uint8_t* randombytes) {
    __shared__ uint8_t smem[MLKEM512Key::shared_memory_size];
    int b = blockIdx.x;
    MLKEM512Key().execute(
        public_keys + b * MLKEM512Key::public_key_size,
        secret_keys + b * MLKEM512Key::secret_key_size,
        randombytes + b * MLKEM512Key::entropy_size,
        workspace   + b * MLKEM512Key::workspace_size, smem);
}

#include <pk.hpp>
using namespace cupqc;

using MLDSA44Sign = decltype(ML_DSA_44()
                 + Function<function::Sign>()
                 + Block() + BlockDim<128>());

__global__ void sign_kernel(uint8_t* signatures,
                            const uint8_t* messages,
                            size_t message_size,
                            const uint8_t* secret_keys,
                            uint8_t* randombytes,
                            uint8_t* workspace) {
    __shared__ uint8_t smem[MLDSA44Sign::shared_memory_size];
    int b = blockIdx.x;
    MLDSA44Sign().execute(signatures, messages, message_size,
                          secret_keys, randombytes, workspace, smem);
}

Inside cuPQC SDK

Public Key
cuPQC-PK

Public-Key Cryptography

ML-KEMML-DSA
Hashing
cuPQC-Hash

Hash functions and Merkle trees

SHA-2SHA-3SHAKEPoseidon2

Release notes

View all release notes →
0.4.1 Library renaming, Poseidon2 KoalaBear field, and multi-block Merkle trees
0.4.0 Poseidon2 hash function and the Merkle Tree API
0.3.0 Introducing the cuPQC SDK, with public-key and hash function libraries

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