FK20 CUDA
fr.cuh
Go to the documentation of this file.
1 // bls12_381: Arithmetic for BLS12-381
2 // Copyright 2022-2023 Dag Arne Osvik
3 // Copyright 2022-2023 Luan Cardoso dos Santos
4 
5 #ifndef FR_CUH
6 #define FR_CUH
7 
8 #include <stdint.h>
9 
24 typedef uint64_t fr_t[4];
25 
30 extern __constant__ fr_t fr_roots[515];
31 
32 extern __device__ __host__ void fr_fromUint64(fr_t &z, const uint64_t *x);
33 extern __device__ void fr_toUint64(const fr_t &x, uint64_t *z);
34 extern __device__ __host__ void fr_cpy(fr_t &z, const fr_t &x);
35 extern __device__ void fr_reduce4(fr_t &z);
36 extern __device__ void fr_neg(fr_t &z);
37 extern __device__ void fr_x2(fr_t &z);
38 extern __device__ void fr_x3(fr_t &z);
39 extern __device__ void fr_x4(fr_t &z);
40 extern __device__ void fr_x8(fr_t &z);
41 extern __device__ void fr_x12(fr_t &z);
42 extern __device__ void fr_add(fr_t &z, const fr_t &x);
43 extern __device__ void fr_sub(fr_t &z, const fr_t &x);
44 extern __device__ void fr_addsub(fr_t &x, fr_t &y);
45 extern __device__ void fr_sqr(fr_t &z);
46 extern __device__ void fr_mul(fr_t &z, const fr_t &x);
47 extern __device__ void fr_inv(fr_t &z);
48 extern __device__ __host__ void fr_zero(fr_t &z);
49 extern __device__ __host__ void fr_one(fr_t &z);
50 
51 extern __device__ bool fr_eq(const fr_t &x, const fr_t &y);
52 extern __device__ bool fr_neq(const fr_t &x, const fr_t &y);
53 extern __device__ bool fr_nonzero(const fr_t &x);
54 extern __device__ bool fr_iszero(const fr_t &x);
55 extern __device__ bool fr_isone(const fr_t &x);
56 
57 extern __device__ void fr_print(const char *s, const fr_t &x);
58 
59 // Device-side FFT functions
60 
61 extern __device__ void fr_fft(fr_t *output, const fr_t *input);
62 extern __device__ void fr_ift(fr_t *output, const fr_t *input);
63 
64 // Kernel wrappers for device-side FFT functions
65 
66 __global__ void fr_fft_wrapper(fr_t *output, const fr_t *input);
67 __global__ void fr_ift_wrapper(fr_t *output, const fr_t *input);
68 __global__ void fr_eq_wrapper(uint8_t *eq, int count, const fr_t *x, const fr_t *y);
69 
70 #endif
71 
72 // vim: ts=4 et sw=4 si
fr_sub
__device__ void fr_sub(fr_t &z, const fr_t &x)
Calculates the difference of two residues modulo p and stores it into z.
Definition: fr_sub.cu:17
fr_x12
__device__ void fr_x12(fr_t &z)
Multiply the residue mod r z by 12 with weak reduction.
Definition: fr_x12.cu:13
fr_x4
__device__ void fr_x4(fr_t &z)
Multiply z by 4, and stores in z, with weak reduction.
Definition: fr_x4.cu:13
fr_fromUint64
__device__ __host__ void fr_fromUint64(fr_t &z, const uint64_t *x)
Converts from uint64_t[4] to a residue modulo r, without reduction.
Definition: fr.cu:59
fr_sqr
__device__ void fr_sqr(fr_t &z)
Squares the value in z as a residue modulo r, and stores back into z.
Definition: fr_sqr.cu:14
fr_x2
__device__ void fr_x2(fr_t &z)
Multiply z by 2, and stores in z, with weak reduction.
Definition: fr_x2.cu:13
fr_print
__device__ void fr_print(const char *s, const fr_t &x)
prints the canonical representation of x to STDOUT.
Definition: fr.cu:41
fr_x8
__device__ void fr_x8(fr_t &z)
Multiply z by 8, and stores in z, with weak reduction.
Definition: fr_x8.cu:13
fr_eq
__device__ bool fr_eq(const fr_t &x, const fr_t &y)
Compares two residues modulo r.
Definition: fr_eq.cu:13
fr_neq
__device__ bool fr_neq(const fr_t &x, const fr_t &y)
Compares two fr_t residues.
Definition: fr_neq.cu:15
fr_t
uint64_t fr_t[4]
Subgroup element stored as a 256-bit array (a 4-element little-endian array of uint64_t)....
Definition: fr.cuh:24
fr_nonzero
__device__ bool fr_nonzero(const fr_t &x)
Check if the reduced input x is different from zero.
Definition: fr_nonzero.cu:12
fr_x3
__device__ void fr_x3(fr_t &z)
Multiply z by 3, and stores in z, with weak reduction.
Definition: fr_x3.cu:12
fr_iszero
__device__ bool fr_iszero(const fr_t &x)
Checks if the residue x modulo f is congruent to zero.
Definition: fr_iszero.cu:13
fr_neg
__device__ void fr_neg(fr_t &z)
Compute an additive inverse of a residue x modulo r. Stores in x. Subtracts x from the highest multip...
Definition: fr_neg.cu:15
fr_eq_wrapper
__global__ void fr_eq_wrapper(uint8_t *eq, int count, const fr_t *x, const fr_t *y)
Checks equality of two arrays of fr_t, element wise, and store in a byte array.
Definition: fr.cu:99
fr_addsub
__device__ void fr_addsub(fr_t &x, fr_t &y)
Computes the sum and the difference of the arguments, storing back into the arguments: (x,...
Definition: fr_addsub.cu:18
fr_roots
__constant__ fr_t fr_roots[515]
Table for the precomputed root-of-unity values.
Definition: fr_roots.cu:17
fr_zero
__device__ __host__ void fr_zero(fr_t &z)
Sets the value of z to zero.
Definition: fr.cu:15
fr_cpy
__device__ __host__ void fr_cpy(fr_t &z, const fr_t &x)
Copy from x into z.
Definition: fr_cpy.cu:14
fr_isone
__device__ bool fr_isone(const fr_t &x)
Checks if the residue x modulo f is congruent to one.
Definition: fr_isone.cu:13
fr_toUint64
__device__ void fr_toUint64(const fr_t &x, uint64_t *z)
Converts fr_t to uint64_t[4].
Definition: fr.cu:76
fr_reduce4
__device__ void fr_reduce4(fr_t &z)
Reduced the value in fr_t to the field modulus.
Definition: fr_reduce4.cu:16
fr_add
__device__ void fr_add(fr_t &z, const fr_t &x)
Computes the sum of two residues x and z modulo r and stores it in z. Device only function.
Definition: fr_add.cu:16
fr_one
__device__ __host__ void fr_one(fr_t &z)
Sets the value of z to one.
Definition: fr.cu:26
fr_mul
__device__ void fr_mul(fr_t &z, const fr_t &x)
Multiply two residues module r z and x, stores back into z.
Definition: fr_mul.cu:13
fr_fft_wrapper
__global__ void fr_fft_wrapper(fr_t *output, const fr_t *input)
wrapper for fr_fft: FFT for fr_t[512]
Definition: fr_fft.cu:316
fr_inv
__device__ void fr_inv(fr_t &z)
Calculates the multiplicative inverse of z, and stores back.
Definition: fr_inv.cu:33
fr_ift_wrapper
__global__ void fr_ift_wrapper(fr_t *output, const fr_t *input)
wrapper for fr_ift: inverse FFT for fr_t[512]
Definition: fr_fft.cu:345
fr_ift
__device__ void fr_ift(fr_t *output, const fr_t *input)
Inverse FFT for fr_t[512].
Definition: fr_fft.cu:170
fr_fft
__device__ void fr_fft(fr_t *output, const fr_t *input)
FFT over Fr.
Definition: fr_fft.cu:26