FK20 CUDA
g1a.cu
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 #include <stdio.h>
6 
7 #include "fp.cuh"
8 #include "fr.cuh"
9 #include "g1.cuh"
10 
21 __device__ __host__ void g1a_fromUint64(g1a_t &a, const uint64_t *x, const uint64_t *y) {
22  fp_fromUint64(a.x, x);
23  fp_fromUint64(a.y, y);
24 }
25 
36 __device__ __host__ void g1a_fromFp(g1a_t &a, const fp_t &x, const fp_t &y) {
37  fp_cpy(a.x, x);
38  fp_cpy(a.y, y);
39 }
40 
48 __device__ void g1a_fromG1p(g1a_t &a, const g1p_t &p) {
49 
50  // uses a.y as temporary storage for the inverse
51 
52  fp_inv(a.y, p.z);
53 
54  fp_mul(a.x, p.x, a.y);
55  fp_mul(a.y, p.y, a.y);
56 }
57 
65 __device__ __host__ void g1a_cpy(g1a_t &a, const g1a_t &b) {
66  fp_cpy(a.x, b.x);
67  fp_cpy(a.y, b.y);
68 }
69 
77 __device__ __host__ void g1a_print(const char *s, const g1a_t &a) {
78 // printf("%s #x%016lx%016lx%016lx%016lx%016lx%016lx #x%016lx%016lx%016lx%016lx%016lx%016lx\n", s, // clisp
79  printf("%s %016lX%016lX%016lX%016lX%016lX%016lX %016lX%016lX%016lX%016lX%016lX%016lX\n", s, // dc
80 // printf("%s 0x%016lx%016lx%016lx%016lx%016lx%016lx 0x%016lx%016lx%016lx%016lx%016lx%016lx\n", s, // python
81  a.x[5], a.x[4], a.x[3], a.x[2], a.x[1], a.x[0],
82  a.y[5], a.y[4], a.y[3], a.y[2], a.y[1], a.y[0]);
83 }
84 
91 __device__ __host__ void g1a_inf(g1a_t &a) {
92  fp_zero(a.x);
93  fp_zero(a.y);
94 };
95 
102 __device__ __host__ void g1a_gen(g1a_t &a) {
103  a.x[5] = 0x17F1D3A73197D794;
104  a.x[4] = 0x2695638C4FA9AC0F;
105  a.x[3] = 0xC3688C4F9774B905;
106  a.x[2] = 0xA14E3A3F171BAC58;
107  a.x[1] = 0x6C55E83FF97A1AEF;
108  a.x[0] = 0xFB3AF00ADB22C6BB;
109 
110  a.y[5] = 0x08B3F481E3AAA0F1;
111  a.y[4] = 0xA09E30ED741D8AE4;
112  a.y[3] = 0xFCF5E095D5D00AF6;
113  a.y[2] = 0x00DB18CB2C04B3ED;
114  a.y[1] = 0xD03CC744A2888AE4;
115  a.y[0] = 0x0CAA232946C5E7E1;
116 };
117 
118 // vim: ts=4 et sw=4 si
__device__ __host__ void fp_zero(fp_t &z)
Sets z to zero.
Definition: fp.cu:15
__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
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
__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
__device__ __host__ void fp_cpy(fp_t &z, const fp_t &x)
Copy from x into z.
Definition: fp_cpy.cu:14
__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
__device__ __host__ void g1a_fromUint64(g1a_t &a, const uint64_t *x, const uint64_t *y)
Converts arrays of uint64_t into a G1 point in affine coordinates. Each array must be uint64_t....
Definition: g1a.cu:21
__device__ __host__ void g1a_fromFp(g1a_t &a, const fp_t &x, const fp_t &y)
Converts Fp values into a point in G1 in affine coordinates. This function does not validate if the c...
Definition: g1a.cu:36
__device__ __host__ void g1a_inf(g1a_t &a)
Set a to the point-at-infinity (0,0)
Definition: g1a.cu:91
__device__ __host__ void g1a_print(const char *s, const g1a_t &a)
Print a standard representation of a, preceded by the user-set string s.
Definition: g1a.cu:77
__device__ void g1a_fromG1p(g1a_t &a, const g1p_t &p)
Converts a point in projective coordinates into affine coordinates.
Definition: g1a.cu:48
__device__ __host__ void g1a_gen(g1a_t &a)
Sets a to the generator point G1 of bls12_381.
Definition: g1a.cu:102
__device__ __host__ void g1a_cpy(g1a_t &a, const g1a_t &b)
Copy from b into a.
Definition: g1a.cu:65
G1 point in affine coordinates.
Definition: g1.cuh:20
fp_t y
Definition: g1.cuh:21
fp_t x
Definition: g1.cuh:21
G1 point in projective coordinates.
Definition: g1.cuh:27
fp_t z
Definition: g1.cuh:28
fp_t x
Definition: g1.cuh:28
fp_t y
Definition: g1.cuh:28