Types#

Descriptor Accessors#

A composed BigInt descriptor BI (for example decltype(BitWidth<256>() + SM<800>() + Thread())) exposes the following compile-time accessors and member types:

static constexpr unsigned int BI::bit_width#

The configured BitWidth, in bits.

static constexpr unsigned int BI::sm_arch#

The configured SM architecture value.

static constexpr unsigned int BI::tpi#

The configured TPI value (1 for Thread() execution).

static constexpr unsigned int BI::num_limbs#

The number of 32-bit limbs in one instance: bit_width / 32.

type BI::error_policy#

The configured on-error policy type—OnErrorNone, OnErrorTrap, or OnErrorPrintTrap—used to report the failures listed under bigint_error. Defaults to OnErrorNone when the descriptor carries no policy operator. See Operators.

type BI::bigint#

The single-width big-integer type for this descriptor. See bigint below.

type BI::bigint_wide#

The double-width (2 * bit_width-bit) big-integer type produced by wide multiplication and squaring, and consumed by Barrett reduction, Montgomery reduction, and double-width division. See bigint_wide below.

type BI::modulus#

montgomery_modulus specialized for this descriptor’s limb count, error policy, and TPI. See montgomery_modulus below.

bigint#

class bigint#

The single-width, fixed-precision big-integer type. Not named directly in user code; obtained as BI::bigint from a composed descriptor BI. All arithmetic, comparison, and modular operations described in Device functions are members of this type. Storage is num_limbs 32-bit limbs (little-endian, limb 0 least significant), distributed across tpi cooperating warp lanes when tpi > 1.

bigint_wide#

class bigint_wide#

The double-width big-integer type holding 2 * num_limbs limbs, split into independent lo and hi halves of type BI::bigint. Produced by mul_wide, operator*, and square; consumed by reduce_barrett, reduce_montgomery, and bigint_wide::div_rem.

bigint lo#

The low num_limbs-limb half.

bigint hi#

The high num_limbs-limb half.

montgomery_modulus#

class montgomery_modulus : public bigint#

A modulus value with its Montgomery reduction constant m' = -M^-1 mod 2^32 precomputed once at construction, for reuse across repeated to_montgomery / from_montgomery / mul_montgomery / reduce_montgomery calls against the same modulus. Publicly derives from bigint, so the whole bigint interface is available on a modulus and one can be passed wherever a bigint is expected. Constructors are not inherited: the three below are the only ways to build one, so there is no scalar constructor as there is for bigint.

The modulus must be odd. Montgomery reduction requires gcd(R, m) = 1, where R = 2^(num_limbs * 32); this holds only when m is odd. Constructing a montgomery_modulus from an even value leaves m' without a valid inverse, so the precomputed constant is meaningless and every subsequent to_montgomery, mul_montgomery, from_montgomery, or reduce_montgomery call against it silently produces incorrect results.

__device__ montgomery_modulus(const bigint &modulus)#

Constructs from an already-loaded modulus value and precomputes m'. modulus must be odd.

__device__ explicit montgomery_modulus(const uint32_t *ptr)#

Loads the modulus from a limb pointer and precomputes m'. The loaded value must be odd.

__device__ montgomery_modulus(const uint32_t *ptr, unsigned int index)#

Loads instance index of a batched modulus array and precomputes m'. The loaded value must be odd.

__device__ uint32_t m_prime() const#

Returns the precomputed Montgomery constant m' = -M^-1 mod 2^32.

bigint_error#

enum class bigint_error#

Status returned by operations that can fail at run time, and passed to the configured on-error policy’s handle() function (see Operators).

enumerator success#

The operation completed successfully.

enumerator divide_by_zero#

div_rem (or a wrapper such as operator/, operator%, or mod) was called with a zero divisor.

enumerator quotient_overflow#

The true quotient does not fit in the supplied quotient width.

enumerator barrett_divide_by_zero#

setup_barrett was invoked with a zero denominator.

enumerator barrett_input_invariant_violated#

reduce_barrett was invoked with a wide numerator num whose high half is not strictly less than the denominator den (num.hi < den does not hold). den must match the denominator most recently passed to setup_barrett.

enumerator inv_mod_not_invertible#

The value and modulus passed to inv_mod are not coprime, so no inverse exists.

enumerator inv_mod_invalid_modulus#

The modulus passed to inv_mod is zero or one.

enumerator inv_mod_zero_input#

The input value passed to inv_mod is zero.

enumerator inv_mod_even_modulus#

Reserved for an internal odd-modulus-only inverse routine. Not returned by inv_mod.