FK20 CUDA
fp.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 FP_CUH
6 #define FP_CUH
7 
8 #include <stdint.h>
9 
14 typedef uint64_t fp_t[6];
15 
16 extern __device__ __host__ void fp_fromUint64(fp_t &z, const uint64_t *x);
17 extern __device__ void fp_toUint64(uint64_t *z, const fp_t &x);
18 extern __device__ __host__ void fp_cpy(fp_t &z, const fp_t &x);
19 extern __device__ void fp_reduce6(fp_t &z);
20 extern __device__ void fp_neg(fp_t &z, const fp_t &x);
21 extern __device__ void fp_x2(fp_t &z, const fp_t &x);
22 extern __device__ void fp_x3(fp_t &z, const fp_t &x);
23 extern __device__ void fp_x4(fp_t &z, const fp_t &x);
24 extern __device__ void fp_x8(fp_t &z, const fp_t &x);
25 extern __device__ void fp_x12(fp_t &z, const fp_t &x);
26 extern __device__ void fp_add(fp_t &z, const fp_t &x, const fp_t &y);
27 extern __device__ void fp_sub(fp_t &z, const fp_t &x, const fp_t &y);
28 extern __device__ void fp_sqr(fp_t &z, const fp_t &x);
29 extern __device__ void fp_mul(fp_t &z, const fp_t &x, const fp_t &y);
30 extern __device__ void fp_mma(fp_t &z, const fp_t &v, const fp_t &w, const fp_t &x, const fp_t &y);
31 extern __device__ void fp_inv(fp_t &z, const fp_t &x);
32 extern __device__ __host__ void fp_zero(fp_t &z);
33 extern __device__ __host__ void fp_one(fp_t &z);
34 
35 extern __device__ bool fp_eq(const fp_t &x, const fp_t &y);
36 extern __device__ bool fp_neq(const fp_t &x, const fp_t &y);
37 extern __device__ bool fp_nonzero(const fp_t &x);
38 extern __device__ bool fp_iszero(const fp_t &x);
39 extern __device__ bool fp_isone(const fp_t &x);
40 
41 extern __device__ void fp_print(const char *s, const fp_t &x);
42 
43 #endif
44 // vim: ts=4 et sw=4 si
fp_toUint64
__device__ void fp_toUint64(uint64_t *z, const fp_t &x)
Converts from residue modulo p (fp_t) to uint64_t[6]. The converted value is in canonical form.
Definition: fp.cu:75
fp_neq
__device__ bool fp_neq(const fp_t &x, const fp_t &y)
Compares two fp_t residues.
Definition: fp_neq.cu:14
fp_neg
__device__ void fp_neg(fp_t &z, const fp_t &x)
Compute an additive inverse of a residue x modulo p. Stores in z. Subtracts x from the highest multip...
Definition: fp_neg.cu:16
fp_mma
__device__ void fp_mma(fp_t &z, const fp_t &v, const fp_t &w, const fp_t &x, const fp_t &y)
Fp multiply-multiply-add. Fast execution of z = (v*w + x*y) mod p The double-wide products are added ...
Definition: fp_mma.cu:20
fp_print
__device__ void fp_print(const char *s, const fp_t &x)
Prints the canonical representation of x to STDOUT.
Definition: fp.cu:39
fp_add
__device__ void fp_add(fp_t &z, const fp_t &x, const fp_t &y)
Computes the sum of two residues x and y modulo p and stores it in z. Device only function.
Definition: fp_add.cu:17
fp_zero
__device__ __host__ void fp_zero(fp_t &z)
Sets z to zero.
Definition: fp.cu:15
fp_iszero
__device__ bool fp_iszero(const fp_t &x)
Checks if the residue x modulo p is congruent to zero.
Definition: fp_iszero.cu:13
fp_x8
__device__ void fp_x8(fp_t &z, const fp_t &x)
Multiplies x by 8 and stores the result into z.
Definition: fp_x8.cu:15
fp_x2
__device__ void fp_x2(fp_t &z, const fp_t &x)
Multiplies x by 2 and stores the result into z.
Definition: fp_x2.cu:15
fp_x12
__device__ void fp_x12(fp_t &z, const fp_t &x)
Multiplies the residue mod p x by 12 and stores the result into z.
Definition: fp_x12.cu:15
fp_x4
__device__ void fp_x4(fp_t &z, const fp_t &x)
Multiplies x by 4 and stores the result into z.
Definition: fp_x4.cu:15
fp_sqr
__device__ void fp_sqr(fp_t &z, const fp_t &x)
Computes the square of the residue x modulo p and stores it in z.
Definition: fp_sqr.cu:16
fp_one
__device__ __host__ void fp_one(fp_t &z)
Sets z to one.
Definition: fp.cu:26
fp_reduce6
__device__ void fp_reduce6(fp_t &z)
Narrow reduction of a residue modulo p, reducing to the canonical representation.
Definition: fp_reduce6.cu:14
fp_t
uint64_t fp_t[6]
Residue modulo p. Any 384-bit representative of each residue is allowed, and stored as a 6-element li...
Definition: fp.cuh:14
fp_nonzero
__device__ bool fp_nonzero(const fp_t &x)
Check if the reduced input x is different from zero.
Definition: fp_nonzero.cu:12
fp_fromUint64
__device__ __host__ void fp_fromUint64(fp_t &z, const uint64_t *x)
Converts uint64_t[6] to fp_t. After this operation, z represents x mod p.
Definition: fp.cu:58
fp_mul
__device__ void fp_mul(fp_t &z, const fp_t &x, const fp_t &y)
Multiplies two Fp residues x and y, stores in z.
Definition: fp_mul.cu:17
fp_cpy
__device__ __host__ void fp_cpy(fp_t &z, const fp_t &x)
Copy from x into z.
Definition: fp_cpy.cu:14
fp_x3
__device__ void fp_x3(fp_t &z, const fp_t &x)
Multiplies x by 3 and stores the result into z.
Definition: fp_x3.cu:15
fp_inv
__device__ void fp_inv(fp_t &z, const fp_t &x)
Calculates the multiplicative inverse of x and stores in z.
Definition: fp_inv.cu:33
fp_isone
__device__ bool fp_isone(const fp_t &x)
Checks if the residue x modulo p is congruent to one.
Definition: fp_isone.cu:13
fp_eq
__device__ bool fp_eq(const fp_t &x, const fp_t &y)
Compares two residues modulo p.
Definition: fp_eq.cu:14
fp_sub
__device__ void fp_sub(fp_t &z, const fp_t &x, const fp_t &y)
Calculates the difference of two residues modulo p and stores it into z.
Definition: fp_sub.cu:16