Skip to content

Tutorial: BigInt Vector Addition

This tutorial walks you through building a 1024-bit big-integer vector addition application using cuPQC-BigInt. You'll learn how to compose a big-integer descriptor, map four cooperating warp lanes onto a single value, and load and store packed little-endian limb buffers from inside your own kernel.

cuPQC-BigInt requires cuPQC SDK 0.6.0 or newer, the release that introduced it. 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 big-integer vector addition example. The Makefile in this directory will compile all examples, including example_bigint_add.

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 the descriptor operators used to build them. The cucheck.hpp header provides the CUDA_CHECK macro used to validate CUDA API calls.

Step 3: Select the Target Architecture

SM<>() selects the architecture the descriptor generates code for. It is optional in a big-integer descriptor and defaults to SM<800> when omitted, which is also what the host compilation pass uses. Deriving it from __CUDA_ARCH__ keeps the descriptor in step with the -arch=native flag in the Makefile:

#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

Define the 1024-bit big-integer type, with four threads cooperating on each integer:

using BI1024 = decltype(BitWidth<1024>() + TPI<4>() + Warp() BIGINT_EXAMPLE_SM());

BitWidth<1024>() sets the width to 1024 bits. The width must be a multiple of 32, and the limb count is BW / 32, so a 1024-bit value is 32 limbs.
TPI<4>() sets the number of threads per instance to four.
Warp() spreads one value across the TPI<4> consecutive warp lanes, which keeps register pressure and carry-propagation cost down for very wide integers. Exactly one of Thread() or Warp() is required; Thread() instead gives a single thread ownership of the whole value and requires TPI to be unset.
BIGINT_EXAMPLE_SM() expands to the + SM<...>() operator chosen in the previous step.

The descriptor exposes the constants used throughout the rest of the program: BI1024::tpi is the threads-per-instance count, BI1024::num_limbs is the number of 32-bit limbs per value, and BI1024::bigint is the value type itself.

Step 5: Create the Addition Kernel

Create a CUDA kernel that adds one element of each input vector:

__global__ void add_kernel(uint32_t* sums,
                           const uint32_t* global_buf_a,
                           const uint32_t* global_buf_b,
                           unsigned int count)
{
    const unsigned int thread_id = blockIdx.x * blockDim.x + threadIdx.x;
    // each group of 4 threads owns one 1024-bit integer
    const unsigned int bigint_index = thread_id / BI1024::tpi;
    if (bigint_index >= count) {
        return;
    }

    // The indexed constructor reads the integer from the packed limb buffer
    const BI1024::bigint a(global_buf_a, bigint_index);
    const BI1024::bigint b(global_buf_b, bigint_index);
    // 1024-bit integer addition (wraps on overflow)
    const auto c = a + b;
    // store the result in the packed output limb buffer
    c.store(sums, bigint_index);
}

Because the descriptor uses Warp() with TPI<4>(), one value is owned by a group of four lanes rather than by a single thread. Dividing the global thread_id by BI1024::tpi gives the index of the value that group is responsible for, so lanes 0-3 work on value 0, lanes 4-7 work on value 1, and so on. Every lane in a group computes the same bigint_index and participates in the same big-integer operations.

Values are arrays of little-endian 32-bit limbs, so limb 0 is least significant. The indexed constructor and store take an instance index, which makes a packed array of independent big integers loadable and storable directly by batch index. Arithmetic is fixed-width and wraps modulo 21024; there is no carry-out beyond the declared width.

Step 6: Implement the Host Code

The host side allocates the packed limb buffers, copies the inputs to the device, launches the kernel, and copies the sums back. Two small values make host-side verification easy; the same layout supports arbitrary 1024-bit input values:

const unsigned int count = 2;
using host_limb_array = std::array<uint32_t, BI1024::num_limbs>;
static_assert(sizeof(host_limb_array) == BI1024::num_limbs * sizeof(uint32_t));
std::vector<host_limb_array> a(count);
std::vector<host_limb_array> b(count);
std::vector<host_limb_array> sums(count);
a[0][0] = 2u;
a[1][0] = 10u;
b[0][0] = 3u;
b[1][0] = 20u;

uint32_t *d_a = nullptr, *d_b = nullptr, *d_sums = nullptr;
const size_t bytes = a.size() * sizeof(host_limb_array);
CUDA_CHECK(cudaMalloc(reinterpret_cast<void**>(&d_a), bytes));
CUDA_CHECK(cudaMalloc(reinterpret_cast<void**>(&d_b), bytes));
CUDA_CHECK(cudaMalloc(reinterpret_cast<void**>(&d_sums), bytes));
CUDA_CHECK(cudaMemcpy(d_a, a.data(), bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_b, b.data(), bytes, cudaMemcpyHostToDevice));

constexpr unsigned int kThreads = 128;
add_kernel<<<(count * BI1024::tpi + kThreads - 1) / kThreads, kThreads>>>(
    d_sums, d_a, d_b, count);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaMemcpy(sums.data(), d_sums, bytes, cudaMemcpyDeviceToHost));

CUDA_CHECK(cudaFree(d_a));
CUDA_CHECK(cudaFree(d_b));
CUDA_CHECK(cudaFree(d_sums));

Each std::array<uint32_t, BI1024::num_limbs> is exactly one packed value, so a std::vector of them is already in the layout the indexed constructor expects, and the static_assert confirms there is no padding.

Launch geometry is the one place where Warp() changes the arithmetic you are used to: the kernel needs count * BI1024::tpi threads rather than count, because every value consumes tpi lanes. Only the least significant limb of each result is inspected, since the inputs were chosen to fit in a single limb:

// ...
std::printf("  2 + 3   = %u (expected 5)\n", sums[0][0]);
std::printf("  10 + 20 = %u (expected 30)\n\n", sums[1][0]);

if (sums[0][0] != 5u || sums[1][0] != 30u) {
    std::fprintf(stderr, "FAIL: unexpected sums.\n");
    return EXIT_FAILURE;
}

Step 7: Build and Run

The Makefile will build all examples in the folder. Run the vector addition example:

make
./example_bigint_add

The Makefile expects the cuPQC SDK at /usr/local/cupqc-sdk, or at a user-specified path set through the CUPQC_SDK_DIR environment variable. cuPQC-BigInt is a static library, and it is linked with link-time optimization through -dlto -lcupqc-bigint.

Step 8: Understanding the Output

Expected output: 

================================================================
BigInt Vector Addition Example
================================================================

This example demonstrates 1024-bit integer addition using cuPQC SDK.
cuPQC-BigInt is a device-side library: every operation is called from
inside your own kernel rather than launched from the host.

Configuration: 1024-bit width, warp execution, 4 threads per instance
Limb layout:   32 little-endian 32-bit limbs per value

========================================
Results
========================================
  2 + 3   = 5 (expected 5)
  10 + 20 = 30 (expected 30)

Example completed successfully.

Customization Tips

Choose Between Thread() and Warp()

A descriptor must specify exactly one execution mode. Thread() gives a single thread ownership of the entire value and requires TPI to be unset, which suits narrower widths. Warp() spreads one value across TPI<T> consecutive warp lanes, keeping register pressure and carry-propagation cost down for very wide integers:

using BI256  = decltype(BitWidth<256>()  + SM<800>() + Thread());
using BI1024 = decltype(BitWidth<1024>() + SM<800>() + TPI<4>() + Warp());

Remember that switching to Warp() changes your launch geometry, since each value now needs tpi threads. Some operations are thread-execution only: division and remainder, scalar conversions such as to_uint32 and to_uint64, limb access with [], and the trapping error policies.

Change the Bit Width

BitWidth<BW>() is the only required operator in a big-integer descriptor. BW must be a multiple of 32, and the limb count is BW / 32. Descriptors also expose a double-width type for operations that need it:

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

cuPQC-BigInt is built for a specific set of widths per TPI. 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, so link errors are the signal to check your width and TPI against the supported list.

Match Operand Widths Explicitly

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. div_rem is the exception and accepts a narrower divisor.

To widen a value explicitly, zero a buffer of the wider limb count, store the narrow value into its low limbs, and construct the wider type from that buffer.

Learn More