Skip to content

Tutorial: BigInt Modular Multiplication

This tutorial walks you through building a 256-bit big-integer application using cuPQC-BigInt. You'll learn how to add two 256-bit values and compute (a * b) mod m inside your own CUDA kernel, with one value per thread.

cuPQC-BigInt requires cuPQC SDK 0.6.0 or newer, the release that introduced the library. cuPQC 0.6 targets compute capabilities 8.0, 8.6, 8.7, 8.9, and 9.0. Unlike the other cuPQC libraries, cuPQC-BigInt is exposed entirely as __device__ functions, so every operation is called from inside your own kernel rather than launched from the host.

Step 1: Project Setup

Clone the cuPQC repository:

git clone https://github.com/NVIDIA/cuPQC.git
cd cuPQC/examples/bigint

This will download all examples including the addition and modular multiplication example. The Makefile in this directory will compile all examples, including example_bigint_addmulmod.

Step 2: Include Required Headers

Start by including the necessary cuPQC SDK headers:

#include <array>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <vector>

#include <bigint.hpp>
#include <cucheck.hpp>

using namespace cupqc;

The bigint.hpp header provides the fixed-width, unsigned multi-precision integer types and their device-side operations. The cucheck.hpp header provides the CUDA_CHECK macro used to wrap CUDA API calls on the host.

Step 3: Select the Target Architecture

SM<>() selects the architecture the descriptor generates code for. It is optional in a bigint descriptor and defaults to SM<800>, which is what the host compilation pass and any unlisted architecture use. Because the Makefile compiles with -arch=native, the example derives the operator from __CUDA_ARCH__ so the descriptor stays in step with whatever architecture is actually being compiled:

#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ == 900)
#define BIGINT_EXAMPLE_SM() + SM<900>()
#elif defined(__CUDA_ARCH__) && (__CUDA_ARCH__ == 890)
#define BIGINT_EXAMPLE_SM() + SM<890>()
#elif defined(__CUDA_ARCH__) && (__CUDA_ARCH__ == 870)
#define BIGINT_EXAMPLE_SM() + SM<870>()
#elif defined(__CUDA_ARCH__) && (__CUDA_ARCH__ == 860)
#define BIGINT_EXAMPLE_SM() + SM<860>()
#elif defined(__CUDA_ARCH__) && (__CUDA_ARCH__ == 800)
#define BIGINT_EXAMPLE_SM() + SM<800>()
#else
#define BIGINT_EXAMPLE_SM()
#endif

The macro stops at 9.0 because that is the highest SM<> specialization the SDK provides, so a newer architecture falls through to the default rather than naming an SM<> that does not exist.

Step 4: Define the Big-Integer Type

Compose the big-integer type from operators, in the same style as the other cuPQC libraries:

// 256-bit integers, one thread per value.
using BI256 = decltype(BitWidth<256>() + Thread() BIGINT_EXAMPLE_SM());

BitWidth<256>() sets the width to 256 bits. BW must be a multiple of 32, and the limb count is BW / 32, so BI256::num_limbs is 8.
Thread() gives a single thread ownership of the whole value. Exactly one of Thread() or Warp() is required; Warp() instead spreads one value across TPI<T> consecutive warp lanes, which keeps register pressure and carry-propagation cost down for very wide integers.
SM<...>() selects the target architecture, supplied here by BIGINT_EXAMPLE_SM().

Only BitWidth is required to form a complete descriptor.

The descriptor exposes two value types:

using bigint      = typename BI256::bigint;       // 256-bit value
using bigint_wide = typename BI256::bigint_wide;  // 512-bit double-width value

Multiplication and squaring produce a double-width result, which is why bigint_wide exists: for a 256-bit type it carries 512 bits. Both operands of a binary operation must be the same type. Nothing is implicitly promoted or truncated, so mixing widths is a compile error rather than a silent conversion.

Step 5: Create the Kernel

Create a CUDA kernel that processes one value per 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);
}

Values are arrays of little-endian 32-bit limbs, so limb 0 is least significant. Constructors and store optionally take an instance index, which makes a packed array of independent big integers loadable and storable directly by batch index:

BI256::bigint a(global_buf_a, index): load instance index from a packed limb buffer
a + b: fixed-width 256-bit addition; the result wraps modulo 2^256
a.mul_mod(b, m): forms the full double-width product of a and b, then reduces it modulo m
sum.store(sums, index): write the value back to instance index of the output buffer

Both results are stored, so the host can verify the sum and the modular product independently.

Step 6: Implement the Host Code

The host allocates packed limb buffers, copies the inputs to the device, launches the kernel, and copies the results back. Small values keep host-side verification easy; the same layout supports arbitrary 256-bit input values:

int main()
{
    // ... banner and configuration printouts ...

    const unsigned int count = 2;
    using host_limb_array = std::array<uint32_t, BI256::num_limbs>;
    static_assert(sizeof(host_limb_array) == BI256::num_limbs * sizeof(uint32_t));
    std::vector<host_limb_array> a(count, host_limb_array{});
    std::vector<host_limb_array> b(count, host_limb_array{});
    std::vector<host_limb_array> m(count, host_limb_array{});
    std::vector<host_limb_array> sums(count);
    std::vector<host_limb_array> products(count);

    a[0][0] = 7u;
    b[0][0] = 5u;
    m[0][0] = 11u;   // (7 + 5) = 12, (7 * 5) mod 11 = 2
    a[1][0] = 40u;
    b[1][0] = 9u;
    m[1][0] = 13u;   // (40 + 9) = 49, (40 * 9) mod 13 = 3

    // ... cudaMalloc and cudaMemcpy of d_a, d_b, d_m, d_sums, d_products ...

    constexpr unsigned int kThreads = 128;
    add_mulmod_kernel<<<(count + kThreads - 1) / kThreads, kThreads>>>(
        d_sums, d_products, d_a, d_b, d_m, count);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaMemcpy(sums.data(), d_sums, bytes, cudaMemcpyDeviceToHost));
    CUDA_CHECK(cudaMemcpy(products.data(), d_products, bytes, cudaMemcpyDeviceToHost));

    // ... cudaFree of all device buffers ...
}

Because each input value fits in limb 0, verification only has to inspect the least significant limb of each result:

    std::printf("  7 + 5            = %u (expected 12)\n", sums[0][0]);
    std::printf("  (7 * 5)  mod 11  = %u (expected 2)\n", products[0][0]);
    std::printf("  40 + 9           = %u (expected 49)\n", sums[1][0]);
    std::printf("  (40 * 9) mod 13  = %u (expected 3)\n\n", products[1][0]);

    if (sums[0][0] != 12u || products[0][0] != 2u ||
        sums[1][0] != 49u || products[1][0] != 3u) {
        std::fprintf(stderr, "FAIL: unexpected sums or products.\n");
        return EXIT_FAILURE;
    }

    std::printf("Example completed successfully.\n");
    return EXIT_SUCCESS;

Step 7: Build and Run

The Makefile will build all examples in the folder. It expects the cuPQC SDK at /usr/local/cupqc-sdk or a user-specified path, which you can set with the CUPQC_SDK_DIR environment variable or by modifying the Makefile. Run the addition and modular multiplication example:

make
./example_bigint_addmulmod

cuPQC-BigInt is a static library built for a specific set of widths per TPI, and it is linked with LTO (-dlto -lcupqc-bigint). A configuration the type system accepts but the shipped library does not provide — an unsupported width, or a Thread()-only operation used on a warp descriptor — compiles cleanly and then fails to link.

Step 8: Understanding the Output

Expected output: 

================================================================
BigInt Addition and Modular Multiplication Example
================================================================

This example demonstrates two of the most common cuPQC-BigInt
operations: fixed-width addition, which wraps modulo 2^256, and
mul_mod, which computes the full double-width product and then
reduces it modulo m.

Configuration: 256-bit width, thread execution (one value per thread)
Limb layout:   8 little-endian 32-bit limbs per value

========================================
Results
========================================
  7 + 5            = 12 (expected 12)
  (7 * 5)  mod 11  = 2 (expected 2)
  40 + 9           = 49 (expected 49)
  (40 * 9) mod 13  = 3 (expected 3)

Example completed successfully.

Customization Tips

Choose the Right Modular Reduction Path

mul_mod performs a full wide multiply followed by a division-based remainder and requires no precomputation, which makes it the right choice for a one-off operation like the one in this example. When you reuse a single modulus across a loop, a precomputed path amortizes its setup cost instead. Barrett precomputes a reciprocal with setup_barrett, then reduce_barrett reduces a double-width value; setup runs once per modulus, and approx and den_clz are opaque state you pass through unchanged:

BI256::bigint approx;
int den_clz;
BI256::bigint::setup_barrett(m, approx, den_clz);   // once per modulus

BI256::bigint_wide num = a.mul_wide(b);
BI256::bigint rem;
BI256::bigint::reduce_barrett(rem, num, m, approx, den_clz);

The wide numerator must satisfy num.hi < den, which holds for a product when both factors are already reduced.

Stay in the Montgomery Domain

For a chain of repeated multiplications against a fixed modulus, convert once with to_montgomery, multiply with mul_montgomery, and convert back at the end with from_montgomery. Remaining in Montgomery form across the operation chain avoids paying a general reduction on every step.

Make Error Handling Strict During Development

Kernels cannot throw exceptions, and only div_rem and inv_mod return a bigint_error enumerator. Operators and functions like /, %, mod, add_mod, mul_mod, and pow_mod have no return channel for error codes, so cuPQC-BigInt embeds the error-handling policy in the compile-time type descriptor. Because the policy is a template parameter, you can debug aggressively during development and then ship with zero runtime cost:

  • OnErrorNone (default): takes no action; the result is unspecified and the kernel continues.
  • OnErrorTrap: calls __trap() on the offending thread, halting the kernel without printing a message.
  • OnErrorPrintTrap: prints the failing block and thread indices along with the error, then traps. Best kept to debugging, since printf from a kernel is expensive.

Both trapping policies require thread execution (Thread()); warp instances use OnErrorNone only. The BigInt vector addition tutorial covers the remaining execution-mode and bit-width rules, including which operations are thread-only and how to widen a value.

Learn More