Device Functions#

Device functions can be accessed from instances of a descriptor NTT.

Standard NTT#

__device__ void NTT::execute(
Precision *data,
const Precision *twiddles,
const Precision p,
)#

Performs an in-place forward NTT on data.

Parameters:

  • data – Shared-memory buffer of length N, populated by NTT::load() or NTT::load_to_mont(); must already be in Montgomery form. Used in place as both input and output. Must not be a global memory pointer. The forward NTT overwrites data with the result, in Montgomery form.

  • twiddles – Precomputed forward twiddle table of length N, in Montgomery form.

  • p – Prime field modulus.

__device__ void NTT::execute(
Precision *data,
const Precision *inv_twiddles,
const Precision p,
const Precision N_inv,
)#

Performs an in-place inverse NTT on data.

Parameters:

  • data – Shared-memory buffer of length N, populated by NTT::load() or NTT::load_to_mont(); must already be in Montgomery form. Used in place as both input and output. Must not be a global memory pointer. The inverse NTT overwrites data with the result, in Montgomery form.

  • inv_twiddles – Precomputed inverse twiddle table of length N, in Montgomery form.

  • p – Prime field modulus.

  • N_inv – Multiplicative inverse of N modulo p (N-1 mod p), applied as the normalisation factor.

__device__ void NTT::make_twiddles(
Precision *twiddles,
const Precision p,
const Precision g,
)#

Writes N forward twiddle factors to the output buffer twiddles for prime p and primitive root g. Call inside a single-thread kernel (<<<1, 1>>>). The twiddle factors are produced in standard (non-Montgomery) form; call NTT::transform_twiddles_to_mont() afterwards to convert them to Montgomery form before use.

Parameters:

  • twiddles – Output buffer of length N; on return it holds the forward twiddle table in standard form.

  • p – Prime field modulus.

  • g – Primitive root modulo p used to generate twiddles.

__device__ void NTT::transform_twiddles_to_mont(
Precision *twiddles,
const Precision p,
)#

Converts a twiddle table from standard form to Montgomery form in place. Unlike NTT::make_twiddles(), this is a blockwise function — call it from a kernel launched with a full block of NTT::BlockDim threads (not <<<1, 1>>>), with all threads in the block calling it collectively. Call this after NTT::make_twiddles() before passing the twiddle table to any execute function.

Parameters:

  • twiddles – Twiddle table of length N to be converted in place; on return holds the twiddle table in Montgomery form.

  • p – Prime field modulus.

Data Transfer#

These functions move polynomial data between global memory and the shared-memory workspace used by NTT::execute(). load and store copy raw values without any domain conversion — use load when the source data is already in Montgomery form, and store when the result should remain in Montgomery form in global memory. Use load_to_mont / store_from_mont to convert explicitly at the boundary.

__device__ void NTT::load(Precision *sdata, const Precision *data)#

Loads N elements from global memory data into the shared-memory workspace sdata without domain conversion. Use when data is already in Montgomery form.

Parameters:

  • sdata – Shared-memory destination buffer of length N.

  • data – Global-memory source buffer of length N.

__device__ void NTT::load_to_mont(
Precision *sdata,
const Precision *data,
const nttConst<Precision> ntt_const = nttConst<Precision>(),
)#

Loads N elements from global memory data into the shared-memory workspace sdata, converting from standard to Montgomery form using ntt_const.

Parameters:

  • sdata – Shared-memory destination buffer of length N.

  • data – Global-memory source buffer of length N in standard (non-Montgomery) form.

  • ntt_const – Montgomery constants for the prime. Construct with nttConst<Precision>(p).

__device__ void NTT::store(Precision *sdata, Precision *data)#

Stores N elements from the shared-memory workspace sdata to global memory data without domain conversion. Use when the result should remain in Montgomery form.

Parameters:

  • sdata – Shared-memory source buffer of length N.

  • data – Global-memory destination buffer of length N.

__device__ void NTT::store_from_mont(
Precision *sdata,
Precision *data,
const nttConst<Precision> ntt_const = nttConst<Precision>(),
)#

Stores N elements from the shared-memory workspace sdata to global memory data, converting from Montgomery to standard form using ntt_const.

Parameters:

  • sdata – Shared-memory source buffer of length N.

  • data – Global-memory destination buffer of length N.

  • ntt_const – Montgomery constants for the prime. Construct with nttConst<Precision>(p).

Threadwise Montgomery Conversion#

These functions convert a single element between standard and Montgomery form. Unlike NTT::load_to_mont() and NTT::store_from_mont(), which cooperate across all threads in the block to convert the full N-element array, these functions operate on one value per calling thread and can be called independently without synchronisation. Use them for pointwise conversions where only individual elements need to change domain.

__device__ Precision NTT::to_mont(
Precision x,
const nttConst<Precision> ntt_const = nttConst<Precision>(),
)#

Converts a single element x from standard form to Montgomery form. Each calling thread converts its own element independently.

Parameters:

  • x – Element in standard (non-Montgomery) form.

  • ntt_const – Montgomery constants for the prime. Construct with nttConst<Precision>(p).

Returns:

x in Montgomery form.

__device__ Precision NTT::from_mont(
Precision x,
const nttConst<Precision> ntt_const = nttConst<Precision>(),
)#

Converts a single element x from Montgomery form to standard form. Each calling thread converts its own element independently.

Parameters:

  • x – Element in Montgomery form.

  • ntt_const – Montgomery constants for the prime. Construct with nttConst<Precision>(p).

Returns:

x in standard (non-Montgomery) form.

Transformed-Domain Arithmetic#

Element-wise modular arithmetic on two NTT-transformed polynomials held in shared memory. Both operands are assumed to be in Montgomery form, as produced by NTT::execute() after a Montgomery-domain load. data1 is updated in place; data2 is read-only.

__device__ void NTT::add(
Precision *data1,
const Precision *data2,
const Precision p,
)#

Computes data1[i] = (data1[i] + data2[i]) mod p for all i in [0, N). Both buffers must be in Montgomery form.

Parameters:

  • data1 – In/out shared-memory buffer of length N in Montgomery form; updated in place.

  • data2 – Read-only shared-memory buffer of length N in Montgomery form.

  • p – Prime field modulus.

__device__ void NTT::mul(
Precision *data1,
const Precision *data2,
const Precision p,
)#

Computes data1[i] = (data1[i] * data2[i]) mod p for all i in [0, N). Both buffers must be in Montgomery form; the result is also in Montgomery form.

Parameters:

  • data1 – In/out shared-memory buffer of length N in Montgomery form; updated in place.

  • data2 – Read-only shared-memory buffer of length N in Montgomery form.

  • p – Prime field modulus.

__device__ void NTT::sub(
Precision *data1,
const Precision *data2,
const Precision p,
)#

Computes data1[i] = (data1[i] - data2[i]) mod p for all i in [0, N). Both buffers must be in Montgomery form.

Parameters:

  • data1 – In/out shared-memory buffer of length N in Montgomery form; updated in place.

  • data2 – Read-only shared-memory buffer of length N in Montgomery form.

  • p – Prime field modulus.

Staged NTT#

(Requires SubSize<M>() in the descriptor.)

__device__ void NTT::stage_1_execute(
Precision *data,
const Precision *twiddles,
const Precision p,
)#

First pass of a two-stage forward NTT. Launch with M blocks; each block processes K = N/M consecutive elements. Requires shared memory of fwd_stage_1_ntt_shared_workspace_size<N, M, Precision>().

Parameters:

  • data – Shared-memory buffer of size K (= N/M), populated by NTT::stage_1_load() or NTT::stage_1_load_to_mont(); must already be in Montgomery form. Used in place as input and output. Must not be a global memory pointer.

  • twiddles – Precomputed forward twiddle table for this pass, in Montgomery form.

  • p – Prime field modulus.

__device__ void NTT::stage_2_execute(
Precision *data,
const Precision *twiddles,
const Precision p,
)#

Second pass of a two-stage forward NTT. Launch with K blocks; each block processes M strided elements. Requires fwd_stage_2_ntt_shared_workspace_size<N, M, Precision>().

Parameters:

  • data – Shared-memory buffer of size M, populated by NTT::stage_2_load() from the stage 1 output, which is always in Montgomery form. Used in place as input and output. Must not be a global memory pointer.

  • twiddles – Precomputed forward twiddle table for this pass, in Montgomery form.

  • p – Prime field modulus.

__device__ void NTT::stage_1_execute(
Precision *data,
const Precision *inv_twiddles,
const Precision p,
)

First pass of a two-stage inverse NTT. Launch with K blocks; each block processes M strided elements. Requires inv_stage_1_ntt_shared_workspace_size<N, M, Precision>().

Parameters:

  • data – Shared-memory buffer of size M, populated by NTT::stage_1_load() or NTT::stage_1_load_to_mont(); must already be in Montgomery form. Used in place as input and output. Must not be a global memory pointer.

  • inv_twiddles – Precomputed inverse twiddle table for this pass, in Montgomery form.

  • p – Prime field modulus.

__device__ void NTT::stage_2_execute(
Precision *data,
const Precision *inv_twiddles,
const Precision p,
const Precision N_inv,
)#

Second pass of a two-stage inverse NTT; applies the 1/N normalisation. Launch with M blocks; each block processes K = N/M consecutive elements. Requires inv_stage_2_ntt_shared_workspace_size<N, M, Precision>().

Parameters:

  • data – Shared-memory buffer of size K (= N/M), populated by NTT::stage_2_load() from the stage 1 output, which is always in Montgomery form. Used in place as input and output. Must not be a global memory pointer.

  • inv_twiddles – Precomputed inverse twiddle table for this pass, in Montgomery form.

  • p – Prime field modulus.

  • N_inv – Multiplicative inverse of N modulo p (N-1 mod p), applied as the normalisation factor.

Staged Data Transfer#

(Requires SubSize<M>() in the descriptor.)

These functions move polynomial data between global memory and the per-stage shared-memory workspace. Direction is inferred from the descriptor: forward and inverse stages use different access patterns (strided vs. contiguous) so the correct layout is selected automatically. The id argument is the current block index (pass blockIdx.x).

Only the boundary transfers have domain-converting variants: stage_1_load_to_mont converts standard-form input to Montgomery form on the way in, and stage_2_store_from_mont converts back to standard form on the way out. stage_1_store and stage_2_load handle the intermediate global-memory buffer between stage 1 and stage 2, which always stays in Montgomery form, so there is no stage_1_store_from_mont or stage_2_load_to_mont — stage_2_load always assumes its source is already in Montgomery form. stage_1_load and stage_2_store likewise copy data without domain conversion — use them when data is already in or should remain in Montgomery form.

__device__ void NTT::stage_1_load(
Precision *sdata,
const Precision *data,
const int id,
)#

Loads the sub-array for stage 1 from global memory data into sdata without domain conversion. Use when data is already in Montgomery form.

  • Forward (M blocks, K = N/M elements each): reads K strided elements data[id + i*M].

  • Inverse (K blocks, M elements each): reads M contiguous elements data[id*M + i].

Parameters:

  • sdata – Shared-memory destination buffer.

  • data – Global-memory source buffer of length N.

  • id – Block index (blockIdx.x).

__device__ void NTT::stage_1_load_to_mont(
Precision *sdata,
const Precision *data,
const int id,
const nttConst<Precision> ntt_const = nttConst<Precision>(),
)#

Like NTT::stage_1_load() but converts from standard to Montgomery form on load using ntt_const.

Parameters:

  • sdata – Shared-memory destination buffer.

  • data – Global-memory source buffer of length N in standard (non-Montgomery) form.

  • id – Block index (blockIdx.x).

  • ntt_const – Montgomery constants for the prime. Construct with nttConst<Precision>(p).

__device__ void NTT::stage_1_store(
Precision *sdata,
Precision *data,
const int id,
)#

Stores the sub-array computed by stage 1 from sdata to global memory data.

  • Forward: writes strided elements back to data[id + i*M].

  • Inverse: writes contiguous elements back to data[id*M + i].

Parameters:

  • sdata – Shared-memory source buffer.

  • data – Global-memory destination buffer of length N.

  • id – Block index (blockIdx.x).

__device__ void NTT::stage_2_load(
Precision *sdata,
const Precision *data,
const int id,
)#

Loads the sub-array for stage 2 from global memory data into sdata without domain conversion. data is assumed to already be in Montgomery form — the intermediate buffer written by stage_1_store after stage 1 is always in Montgomery form, so there is no stage_2_load_to_mont variant.

  • Forward (K blocks, M elements each): reads M contiguous elements data[id*M + i].

  • Inverse (M blocks, K elements each): reads K strided elements data[id + i*M].

Parameters:

  • sdata – Shared-memory destination buffer.

  • data – Global-memory source buffer of length N.

  • id – Block index (blockIdx.x).

__device__ void NTT::stage_2_store(
Precision *sdata,
Precision *data,
const int id,
)#

Stores the sub-array computed by stage 2 from sdata to global memory data without domain conversion. Use when the result should remain in Montgomery form.

  • Forward: writes contiguous elements back to data[id*M + i].

  • Inverse: writes strided elements back to data[id + i*M].

Parameters:

  • sdata – Shared-memory source buffer.

  • data – Global-memory destination buffer of length N.

  • id – Block index (blockIdx.x).

__device__ void NTT::stage_2_store_from_mont(
Precision *sdata,
Precision *data,
const int id,
const nttConst<Precision> ntt_const = nttConst<Precision>(),
)#

Like NTT::stage_2_store() but converts from Montgomery to standard form on store using ntt_const.

Parameters:

  • sdata – Shared-memory source buffer.

  • data – Global-memory destination buffer of length N.

  • id – Block index (blockIdx.x).

  • ntt_const – Montgomery constants for the prime. Construct with nttConst<Precision>(p).