cuPQC-BigInt: Multi-Precision Integer Arithmetic

The cuPQC-BigInt library provides GPU-optimized, fixed-width, unsigned multi-precision (big) integer arithmetic designed for direct integration into CUDA kernels. Its API is exposed entirely as __device__ functions, so every operation is called directly from within your kernel rather than launched or dispatched from the host. It is a core building block for classical public-key primitives, modular exponentiation, and any workload that needs integers wider than the native 32/64-bit hardware types. As a device-side library, cuPQC-BigInt integrates directly into your CUDA kernels, enabling you to fuse big-integer arithmetic with other GPU computations for maximum performance.

Execution Modes

The cuPQC-BigInt library provides two execution modes that determine how a single big-integer instance is distributed across CUDA threads:

• Thread execution: A single CUDA thread owns and operates on the entire big integer. Selected with the Thread() operator. Requires TPI to be unset (implicitly 1).

• Warp execution: TPI (Threads Per Instance) consecutive lanes of a warp cooperate on a single big-integer instance, with each lane holding a contiguous slice of the limbs. Selected with the Warp() operator together with TPI<T>. Multiple independent instances can be packed into one warp—for example TPI<8> packs four independent instances into a 32-lane warp, each computed cooperatively by 8 lanes.

Exactly one of Thread() or Warp() must be specified in a descriptor.

The notation in the panel below is used throughout the rest of this page for widths, lane counts, and modular arithmetic.

Notation for BW, T, and m

Bit width BW is a multiple of 32, set with BitWidth<BW>; the limb count is BW / 32.

Threads per instance T is a power of two, 1 <= T <= 32, set with TPI<T>. It may not exceed the limb count.

m is the modulus of a modular operation, and a, b and e are its operands and exponent.

Bit Width and Limb Layout

Every big-integer value is a fixed-width unsigned integer made of 32-bit limbs:

Aspect	Constraint	Detail
Bit width BitWidth<BW>()	BW = 32 * k, k > 0	Must be a multiple of 32 and greater than 0. The number of limbs is BW / 32.
Limb order	Little endian	Limb 0 is least significant. Values are stored and loaded as arrays, plain or batched, of uint32_t.
Threads per instance TPI<T>()	1 · 2 · 4 · 8 · 16 · 32 Also no larger than the limb count BW / 32.	Required with Warp(), forbidden with Thread(); every value in this range is validated. Limbs are distributed across the TPI cooperating lanes with the lowest-index lane holding the least-significant limbs; any padding limbs are placed at the top of the highest-index lane.
Wide results bigint_wide	2 * BW bits	Split into lo and hi halves. Produced by mul_wide, operator*, and square, and consumed by reduce_barrett, reduce_montgomery, and bigint_wide::div_rem.

Supported Widths

BitWidth accepts any multiple of 32, but the shipped library is built for a specific set of widths per TPI. The ranges below are inclusive, and every integer limb count inside them is built rather than only the powers of two, so 37 limbs (1184 bits) at TPI<4> is a real configuration.

TPI	Limb counts	Bit widths
1	1 · 2 · 4 · 8–32	32 · 64 · 128 · 256–1024
2	4–64	128–2048
4	8–128	256–4096
8	16–128	512–4096
16	32–128	1024–4096
32	64–128	2048–4096

Two rules shape the table. Each warp configuration starts at two limbs per lane, 2 * T, and stops at 32 limbs per lane or 128 limbs in total, whichever is smaller. That is why thread execution ends at 32 limbs and TPI<2> at 64, while everything from TPI<4> up ends at 128.

Thread execution is the exception at the bottom, adding 1, 2, and 4 limbs below its range: 32-, 64-, and 128-bit thread instances work, while 3, 5, 6, and 7 limbs are gaps. A width outside these sets satisfies the type system and then fails to link.

Operand Widths

Every binary operation requires both operands to be the same type: the same bit width, the same threads per instance, and the same error policy. Nothing is implicitly promoted, truncated, or converted between widths, so passing a 256-bit value where a 512-bit one is expected is a compile error rather than a silent conversion. Adding two BW-bit values yields a BW-bit result; multiplying them yields a 2 * BW-bit bigint_wide.

div_rem is the one operation that accepts operands of differing widths:

Overload	Dividend	Divisor and remainder	Quotient
bigint::div_rem	BW	BW / 2	BW or BW / 2
bigint_wide::div_rem	2 * BW	BW	BW

The remainder always matches the divisor. A quotient that is neither the dividend width nor half of it is rejected at compile time, and the divisor widths above are the ones the shipped library provides—other combinations the type system would accept are not built, so they fail to link rather than to compile.

Everywhere else, a narrower value has to be widened explicitly first. There is no in-register widening: zero a buffer of the wider limb count, store the narrow value into its low limbs, and construct the wider type from that buffer.

using BI256 = decltype(BitWidth<256>() + SM<800>() + Thread());
using BI512 = decltype(BitWidth<512>() + SM<800>() + Thread());

uint32_t buf[512 / 32] = {};        // zero-filled
narrow.store(buf);                  // 256-bit value occupies the low limbs
typename BI512::bigint wide(buf);   // same value, 512-bit type

Scalar operands are a separate facility and do not imply mixed-width support. Addition, subtraction, and multiplication take a uint32_t; comparisons take a uint32_t or uint64_t on either side; pow_mod takes a uint64_t exponent; a bigint wider than one limb can be constructed from a uint64_t; and to_uint32 and to_uint64 read a scalar back out. There are no uint64_t overloads for addition, subtraction, or multiplication.

Supported Operations

Each function name below links to its entry in the Device functions reference; the section and category labels link to the matching headings there.

Category		Functions	Notes
Construction and Access		bigint store to_uint32 to_uint64 []	bigint — construct from a uint32_t/uint64_t scalar, or from a limb pointer (plain or batched). store — write back to global memory the same way. to_uint32 / to_uint64 — read the low bits back as a scalar. [] — index individual limbs in thread execution.
Basic Arithmetic	Addition and Subtraction	add_scalar sub_scalar + - += -=	Fixed-width, bigint-bigint and bigint-scalar. Results wrap modulo 2BW.
	Multiplication	mul_scalar mul_low mul_high mul_wide square * *=	mul_low — about half the cost of a full wide multiply. mul_wide / square — double-width product. * — full double-width product; *= multiplies in place by a scalar.
	Division	div_rem /	Same-width divisor, quotient and remainder — or independently sized divisor and quotient widths. The bigint_wide overload divides a double-width dividend.
Comparison		compare == != < <= > >=	Three-way compare against another instance, a uint32_t, or a uint64_t. Full relational and equality set, including scalar forms on either side.
Bitwise and Shift		clz & | ^ ~ &= |= ^= << >> <<= >>=	Bitwise and logical shift operators over the full width, plus a leading-zero count.
Modular Arithmetic	Direct	mod add_mod sub_mod mul_mod pow_mod inv_mod %	mod, % — plain remainder. add_mod, sub_mod — operands must already be reduced. mul_mod — full multiply plus remainder, not Montgomery. pow_mod — uint64_t or bigint exponent. inv_mod — odd and even moduli, scalar and warp-cooperative.
	Barrett	setup_barrett reduce_barrett	Static helpers for fast remainder across repeated reductions against the same modulus.
	Montgomery	to_montgomery from_montgomery mul_montgomery reduce_montgomery	REDC-based conversion and Montgomery multiplication. montgomery_modulus precomputes the constant m' = -m-1 mod 232 once, and reduce_montgomery reduces a double-width value.

Warp Execution Coverage

Warp execution (TPI > 1) covers most of the interface, including all modular arithmetic, Barrett, and Montgomery. The following are built for Thread() execution only:

• div_rem, operator/, and operator%, including bigint_wide::div_rem. Use mod where only a remainder is needed.

• to_uint32 and to_uint64.

• The uint64_t constructor.

• operator+= and operator-= with a bigint right-hand operand; the uint32_t forms are available.

• operator[], which no single lane can satisfy when the instance spans several.

• OnErrorTrap and OnErrorPrintTrap; OnErrorNone is the only policy available.

Calling one of these on a warp descriptor compiles and then fails to link.

Choosing Modular Reduction

cuPQC-BigInt exposes several ways to reduce modulo a user modulus. They differ mainly in whether they pay a division on every operation or precompute once to avoid it, so the paths below are grouped by how much setup each one needs:

Setup	Operation	Functions	Notes
None One-off use	Add / subtract (a ± b) mod m	add_mod, sub_mod	Conditional reduction, no division. Both operands must already be reduced (a < m, b < m).
	Remainder a mod m	mod, %, div_rem	Plain division. Prefer div_rem when you need a status code or the quotient.
	Multiply (a * b) mod m	mul_mod	Wide multiply, then a division to reduce.
Precompute Many operations, same m	Reduce wide values any m	setup_barrett, reduce_barrett	Barrett. Trades the per-call division for multiplies, so it pays off once you reduce by the same m repeatedly. Takes the double-width results of mul_wide and square.
	Chain multiplications odd m only	to_montgomery, mul_montgomery, reduce_montgomery, from_montgomery	Montgomery. Also division-free, and the more multiplications you chain the further the one-time conversion is amortized. Convert once, stay in Montgomery form, convert back at the end.
Internal Chosen for you	Exponentiate ae mod m	pow_mod	Picks one of the above for you: odd m uses Montgomery internally, even m uses mul_mod. Not constant-time

Error Handling

Operations that can fail at run time—division, remainder, modular reduction, and modular inverse—report failures through a bigint_error enumerator (success, divide_by_zero, quotient_overflow, barrett_divide_by_zero, barrett_input_invariant_violated, inv_mod_not_invertible, inv_mod_invalid_modulus, inv_mod_zero_input, inv_mod_even_modulus).

Every failure is additionally routed through a compile-time on-error policy, selected by adding one of the following operator types to the descriptor:

Policy	On failure	Use when
OnErrorNone (default)	Nothing. The result is unspecified and the kernel continues.	You inspect the returned bigint_error (from div_rem / inv_mod) yourself, or the inputs are already known to be valid and you want no added cost.
OnErrorTrap	Calls __trap() on the offending thread, halting the kernel. No message is printed.	Catching errors from operations that do not themselves return a status, such as operator/, operator%, mod, add_mod, sub_mod, mul_mod, and pow_mod.
OnErrorPrintTrap	Prints the error together with the offending block and thread indices, then traps.	You are debugging. Prefer one of the above in production, since printf from a kernel is expensive.

Only div_rem and inv_mod return a bigint_error, so under OnErrorNone they are the only operations whose failures a caller can detect. For the rest—operator/, operator%, mod, add_mod, sub_mod, mul_mod, pow_mod, to_montgomery, setup_barrett and reduce_barrett—the policy is the only thing that reports a failure, and the default policy reports nothing. A trapping policy is what makes those failures observable.

Key Features

• High Performance: GPU-optimized limb kernels exploit PTX-level carry-chain and multiply-add intrinsics for maximum throughput.

• Predictable Width Support: Each TPI builds a specific set of widths from 32 to 4096 bits; see Supported Widths.

• Choosing an Execution Model: Single-thread ownership suits small and medium widths; warp-cooperative execution (TPI) spreads very wide integers—and their carry propagation—across multiple lanes.

• Montgomery Domain Support: to_montgomery/from_montgomery/mul_montgomery and a precomputed-constant montgomery_modulus type make REDC-based modular multiplication efficient across repeated operations against the same modulus.

• Configurable Error Handling: Choose whether arithmetic failures are silent, trap, or trap-with-diagnostics by adding OnErrorTrap() or OnErrorPrintTrap() to the descriptor.

• Batch-Friendly Storage: Constructors and store accept an instance index, so an array of independent big integers can be loaded and stored directly by batch index.