FK20 CUDA
g1p_sub.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 "g1.cuh"
9 
17 __device__ void g1p_sub(g1p_t &p, const g1p_t &q) {
18 
19  fp_t
20  X1, Y1, Z1,
21  X2, Y2, Z2,
22  t0, t1, t2, t3;
23 
24  fp_cpy(X1, p.x);
25  fp_cpy(Y1, p.y);
26  fp_cpy(Z1, p.z);
27 
28  fp_cpy(X2, q.x);
29  fp_cpy(Y2, q.y);
30  fp_cpy(Z2, q.z);
31 
32  if (!g1p_isInf(q))
33  fp_neg(Y2, Y2);
34 
35  // Adapted from eprint 2015-1060, algorithm 7.
36  // Modified to avoid overwriting inputs and remove one temp value.
37  // 12 mul, 0 square, 11 add, 5 sub, 2 x12, 1 x3.
38 
39  fp_add(t0, X1, Y1); // t3
40 
41  //fp_print("T3 = ", t0);
42 
43  fp_add(t1, Y1, Z1); // t8
44 
45  //fp_print("T8 = ", t1);
46 
47  fp_add(t2, Z1, X1); // td
48 
49  //fp_print("Td = ", t2);
50 
51 
52  fp_mul(X1, X1, X2); // t0
53 
54  //fp_print("T0 = ", X1);
55 
56  fp_mul(Y1, Y1, Y2); // t1
57 
58  //fp_print("T1 = ", Y1);
59 
60  fp_mul(Z1, Z1, Z2); // t2
61 
62  //fp_print("T2 = ", Z1);
63 
64 
65  fp_add(t3, X2, Y2); // t4
66 
67  //fp_print("T4 = ", t3);
68 
69  fp_add(Y2, Y2, Z2); // t9
70 
71  //fp_print("T9 = ", Y2);
72 
73  fp_add(Z2, Z2, X2); // te
74 
75  //fp_print("Te = ", Z2);
76 
77 
78  fp_mul(X2, t3, t0); // t5
79 
80  //fp_print("T5 = ", X2);
81 
82  fp_mul(Y2, Y2, t1); // ta
83 
84  //fp_print("Ta = ", Y2);
85 
86  fp_mul(Z2, Z2, t2); // tf
87 
88  //fp_print("Tf = ", Z2);
89 
90 
91  fp_x3(t0, X1); // ti
92 
93  //fp_print("Ti = ", t0);
94 
95  fp_add(t1, Y1, Z1); // tb
96 
97  //fp_print("Tb = ", t1);
98 
99  fp_add(t2, Z1, X1); // tg
100 
101  //fp_print("Tg = ", t2);
102 
103  fp_x12(t3, Z1); // tk
104 
105  //fp_print("Tk = ", t3);
106 
107 
108  fp_add(X1, X1, Y1); // t6
109 
110  //fp_print("T6 = ", X1);
111 
112  fp_add(Z1, Y1, t3); // tl
113 
114  //fp_print("Tl = ", Z1);
115 
116  fp_sub(Y1, Y1, t3); // tm
117 
118  //fp_print("Tm = ", Y1);
119 
120 
121  fp_sub(X1, X2, X1); // t7
122 
123  //fp_print("T7 = ", X1);
124 
125  fp_mul(X2, X1, t0); // ts
126 
127  //fp_print("Ts = ", X2);
128 
129 
130  fp_mul(X1, X1, Y1); // tp
131 
132  //fp_print("Tp = ", X1);
133 
134  fp_mul(Y1, Y1, Z1); // tr
135 
136  //fp_print("Tr = ", Y1);
137 
138 
139  fp_sub(Y2, Y2, t1); // tc
140 
141  //fp_print("Tc = ", Y2);
142 
143  fp_mul(Z1, Z1, Y2); // tt
144 
145  //fp_print("Tt = ", Z1);
146 
147  fp_sub(Z2, Z2, t2); // th
148 
149  //fp_print("Th = ", Z2);
150 
151 
152  fp_x12(Z2, Z2); // tn
153 
154  //fp_print("Tn = ", Z2);
155 
156  fp_mul(Y2, Y2, Z2); // to
157 
158  //fp_print("To = ", Y2);
159 
160  fp_mul(Z2, Z2, t0); // tq
161 
162  //fp_print("Tq = ", Z2);
163 
164 
165  fp_sub(X1, X1, Y2); // X3
166  fp_add(Y1, Y1, Z2); // Y3
167  fp_add(Z1, Z1, X2); // Z3
168 
169  //fp_print("X4 = ", X1);
170  //fp_print("Y4 = ", Y1);
171  //fp_print("Z4 = ", Z1);
172 
173  fp_cpy(p.x, X1);
174  fp_cpy(p.y, Y1);
175  fp_cpy(p.z, Z1);
176 }
177 
178 // vim: ts=4 et sw=4 si
__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
__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
__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
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_x3(fp_t &z, const fp_t &x)
Multiplies x by 3 and stores the result into z.
Definition: fp_x3.cu:15
__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
__device__ bool g1p_isInf(const g1p_t &p)
Check if the value stored in p is the the/any point at infinity. This implementation uses (0,...
Definition: g1p_ispoint.cu:20
__device__ void g1p_sub(g1p_t &p, const g1p_t &q)
Point subtraction using projective coordinates. p ← p-q.
Definition: g1p_sub.cu:17
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