cuPQC-BigInt Usage#
The cuPQC-BigInt library provides GPU-optimized, fixed-width multi-precision integer
arithmetic that integrates directly into CUDA kernels: every operation is exposed as a
__device__ function, designed to be called directly from within your kernel rather than
launched or dispatched from the host. A BigInt descriptor is defined by
combining cuPQC-BigInt operators to specify the bit
width, execution mode, and (for warp execution) the number of cooperating lanes at compile
time.
Defining a BigInt Descriptor#
A descriptor specifies all instance parameters as a C++ type using operator composition, then
exposes a nested bigint type (and a double-width bigint_wide type) to use in device code.
Thread execution — one thread owns the whole value. Use this for small to moderately wide integers:
#include <bigint.hpp>
using namespace cupqc;
using BI256 = decltype(BitWidth<256>() + SM<800>() + Thread());
using bigint = typename BI256::bigint; // 256-bit value, single thread
using bigint_wide = typename BI256::bigint_wide; // 512-bit (double-width) value
Warp execution — TPI (Threads Per Instance) consecutive lanes cooperate on one value.
Use this for very wide integers (RSA-class and beyond), where spreading limbs and carry
propagation across lanes keeps per-thread register pressure and instruction count down:
using BI2048 = decltype(BitWidth<2048>() + SM<800>() + TPI<32>() + Warp());
using bigint = typename BI2048::bigint; // 2048-bit value, 32 cooperating lanes
using bigint_wide = typename BI2048::bigint_wide; // 4096-bit (double-width) value
TPI must be a power of two between 1 and 32, and must not exceed the instance’s limb
count (BitWidth / 32). A warp can host multiple independent instances at once: with
TPI<8>, a 32-lane warp computes four independent 2048/8=256-bit-per-lane instances in
parallel, one per group of 8 lanes.
Constructing and Storing Values#
Construct a value from a scalar, from a limb pointer, or from a batched array using an
instance index; store mirrors each of the pointer-based constructors:
__global__ void construct_examples_kernel(uint32_t* out, const uint32_t* a, const uint32_t* batch)
{
using bigint = typename BI256::bigint;
// One bigint instance per TPI-lane group, not per block: this launches
// (blockDim.x / BI256::tpi) independent instances per block instead of just one.
const unsigned int global_lane = blockIdx.x * blockDim.x + threadIdx.x;
const unsigned int index = global_lane / BI256::tpi;
bigint from_scalar(7u); // value 7
bigint from_ptr(a); // BIT_WIDTH/32 limbs starting at `a`
bigint from_batch(batch, index); // instance `index` of a batched array
from_scalar.store(out); // store starting at `out`
from_scalar.store(out, index); // store to instance `index` of a batched array
}
For warp-executed instances, every cooperating lane in the group must call the same
constructor/store with the same pointer and index (as computed above, index is the same
for all BI256::tpi lanes in a group); each lane reads and writes only the limbs it owns.
Arithmetic, Bitwise, and Comparison Operations#
Big integers support the familiar arithmetic, bitwise, shift, and comparison operators,
plus scalar (uint32_t) overloads for the common cases:
__global__ void arithmetic_kernel(uint32_t* out, const uint32_t* a, const uint32_t* b)
{
using bigint = typename BI256::bigint;
bigint av(a), bv(b);
bigint sum = av + bv;
bigint diff = av - bv;
bigint shifted = av << 4;
bigint masked = av & bv;
bool is_less = av < bv;
int ordering = av.compare(bv); // <0, 0, or >0
sum.store(out);
}
Multiplication and Wide Results#
A full product of two BW-bit values is 2 * BW bits wide, so multiplication returns a
double-width bigint_wide (with lo/hi halves) unless you only need half the result:
__global__ void multiply_kernel(uint32_t* product_out, uint32_t* squared_out,
const uint32_t* a, const uint32_t* b)
{
using bigint = typename BI256::bigint;
bigint av(a), bv(b);
bigint low_half = av.mul_low(bv); // low BW bits; half the cost of a full multiply
bigint high_half = av.mul_high(bv); // high BW bits
auto product = av.mul_wide(bv); // full 2*BW-bit product; product.lo / product.hi
auto squared = av.square(); // 2*BW-bit square
// bigint_wide's own store() writes both halves (2 * num_limbs limbs) in one call,
// rather than storing product.lo / product.hi separately.
product.store(product_out);
squared.store(squared_out);
}
Division and Remainder#
div_rem returns a bigint_error status and writes the quotient and remainder
by reference. The divisor, quotient, and remainder can all be the same width as the dividend,
or the quotient and divisor can be independently sized:
__global__ void div_rem_kernel(cupqc::bigint_error* status, uint32_t* q, uint32_t* r,
const uint32_t* a, const uint32_t* b)
{
using bigint = typename BI256::bigint;
bigint av(a), bv(b), qv, rv;
*status = av.div_rem(bv, qv, rv); // qv = a / b, rv = a % b
qv.store(q);
rv.store(r);
}
// Convenience wrappers route failures through the OnError policy instead of
// returning a status:
__global__ void div_rem_operators_kernel(uint32_t* q, uint32_t* r,
const uint32_t* a, const uint32_t* b)
{
using bigint = typename BI256::bigint;
bigint av(a), bv(b);
bigint quotient = av / bv;
bigint remainder = av % bv;
quotient.store(q);
remainder.store(r);
}
For a quotient narrower or wider than the divisor, pass differently sized bigint types explicitly:
using narrow_bigint = typename decltype(BitWidth<128>() + SM<800>() + Thread())::bigint;
using wide_bigint = typename BI256::bigint; // 256-bit divisor/remainder
__global__ void div_rem_mixed_kernel(uint32_t* q, uint32_t* r,
const uint32_t* a, const uint32_t* b)
{
wide_bigint av(a), bv(b), rv;
narrow_bigint qv;
av.div_rem(bv, qv, rv); // qv is 128 bits, bv/rv are 256 bits
qv.store(q);
rv.store(r);
}
Modular Arithmetic#
For choosing among *_mod, Barrett, and Montgomery, see
Choosing Modular Reduction on the Features page.
add_mod, sub_mod, and mul_mod combine an arithmetic step with a reduction against a
supplied modulus (mul_mod performs a full multiply followed by remainder, not Montgomery
multiplication). pow_mod accepts either a uint64_t or a bigint exponent:
__global__ void mod_arithmetic_kernel(uint32_t* out, const uint32_t* a, const uint32_t* b,
const uint32_t* m, const uint32_t* e)
{
using bigint = typename BI256::bigint;
bigint av(a), bv(b), mv(m), ev(e);
bigint sum_mod = av.add_mod(bv, mv);
bigint prod_mod = av.mul_mod(bv, mv);
bigint pow_u64 = av.pow_mod(uint64_t{65537}, mv); // e.g. RSA public exponent
bigint pow_big = av.pow_mod(ev, mv);
pow_big.store(out);
}
Modular Inverse#
inv_mod computes the multiplicative inverse of a value modulo m, supporting both odd
and even moduli, and both single-thread and warp-cooperative instances. It returns a
bigint_error describing why the inverse does not exist (if it does not):
__global__ void inv_mod_kernel(cupqc::bigint_error* status, uint32_t* out,
const uint32_t* a, const uint32_t* m)
{
using bigint = typename BI256::bigint;
bigint av(a), mv(m), inverse;
*status = av.inv_mod(mv, inverse); // success, or an inv_mod_* error
inverse.store(out);
}
Montgomery Arithmetic#
For workloads that repeat many modular multiplications against the same modulus (such as
pow_mod’s square-and-multiply ladder), convert operands into the Montgomery domain once,
multiply with mul_montgomery (which folds in the REDC reduction), and convert back with
from_montgomery. Build a montgomery_modulus once per modulus; it precomputes the
Montgomery constant m' = -M^-1 mod 2^32 so it doesn’t need to be recomputed on every call:
__global__ void montgomery_mul_kernel(uint32_t* out, const uint32_t* a, const uint32_t* b,
const uint32_t* m)
{
using bigint = typename BI256::bigint;
using modulus_t = typename BI256::modulus; // montgomery_modulus<...>
bigint av(a), bv(b);
modulus_t mv(m); // precomputes m' once
bigint aR = av.to_montgomery(mv); // a * R mod m
bigint bR = bv.to_montgomery(mv); // b * R mod m
bigint abR = aR.mul_montgomery(bR, mv); // a * b * R mod m
bigint prod = abR.from_montgomery(mv); // a * b mod m
prod.store(out);
}
bigint_wide::reduce_montgomery performs the Montgomery reduction of an already double-width
value (for example the output of mul_wide) directly against a montgomery_modulus,
without a separate multiply step:
__global__ void wide_reduce_kernel(uint32_t* out, const uint32_t* wide_value,
const uint32_t* m)
{
using bigint_wide = typename BI256::bigint_wide;
using modulus_t = typename BI256::modulus;
bigint_wide value(wide_value);
modulus_t mv(m);
value.reduce_montgomery(mv).store(out); // (value * R^-1) mod m
}
Error Handling#
By default (OnErrorNone), failures are only reported through the returned
bigint_error from div_rem and inv_mod; other operations that can fail
(operator/, operator%, mod, add_mod, sub_mod, mul_mod, pow_mod) do
not return a status. Add OnErrorTrap() or OnErrorPrintTrap() to the descriptor to have
every failure trap the kernel, with the latter also printing a diagnostic message identifying
the failing block and thread:
using BI256Checked = decltype(BitWidth<256>() + SM<800>() + Thread()
+ OnErrorPrintTrap());
using bigint = typename BI256Checked::bigint;
__global__ void checked_div_kernel(uint32_t* out, const uint32_t* a, const uint32_t* b)
{
bigint av(a), bv(b);
bigint quotient = av / bv; // prints a diagnostic and traps on divide-by-zero
quotient.store(out);
}
Example: Addition and Modular Multiplication#
The following minimal, single-thread example shows two of the most common cuPQC-BigInt
operations end to end: plain addition and modular multiplication, using a 256-bit
Thread descriptor:
using BI256 = decltype(BitWidth<256>() + SM<800>() + Thread());
__global__ void add_mulmod_kernel(uint32_t* sums, uint32_t* products,
const uint32_t* global_buf_a,
const uint32_t* global_buf_b,
const uint32_t* global_buf_m,
unsigned int count)
{
const unsigned int index = blockIdx.x * blockDim.x + threadIdx.x;
if (index >= count) {
return;
}
const BI256::bigint a(global_buf_a, index);
const BI256::bigint b(global_buf_b, index);
const BI256::bigint m(global_buf_m, index);
// 256-bit integer addition (wraps on overflow).
const auto sum = a + b;
// (a * b) mod m.
const auto product = a.mul_mod(b, m);
sum.store(sums, index);
product.store(products, index);
}
See examples/example_bigint_addmulmod.cu in the SDK for the complete, self-contained
example, including host-side setup and verification.
Compilation#
Include bigint.hpp and link with the cuPQC-BigInt static library using LTO flags:
nvcc -std=c++17 -dlto -arch=sm_80 \
-I${PATH_TO_CUPQC_SDK_INCLUDE} \
-L${PATH_TO_STATIC_LIB} -lcupqc-bigint \
my_bigint_program.cu -o my_bigint_program
For CMake-based projects:
target_link_libraries(YourProgram PRIVATE cupqc-bigint_static)
For detailed installation and compilation instructions, see the Getting Started guide.
API Reference:
cuPQC-BigInt API Reference - cuPQC-BigInt API documentation