FK20 CUDA
g1test_kat.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 "g1.cuh"
8 #include "g1test.cuh"
9 
10 __managed__ g1p_t
11  g1p_x0 = {
12  { 0, 0, 0, 0, 0, 0 },
13 #if G1P_ANYINF
14  { 1, 2, 3, 4, 5, 6 },
15 #else
16  { 1, 0, 0, 0, 0, 0 },
17 #endif
18  { 0, 0, 0, 0, 0, 0 },
19  },
20  g1p_x2 = {
21  { 0x8a13c2b29f4325ad, 0x7803e16723f9f147, 0x6e7f23d4350c60bd, 0x062e48a6104cc52f, 0x9b6d4dac3f33e92c, 0x05dff4ac6726c6cb, }, // 0x05dff4ac6726c6cb9b6d4dac3f33e92c062e48a6104cc52f6e7f23d4350c60bd7803e16723f9f1478a13c2b29f4325ad
22  { 0x2aa4d8e1b1df7ca5, 0x599a8782a1bea48a, 0x395952af0b6a0dbd, 0xd6093a00bb3e2bc9, 0x3c604c0410e5fc01, 0x14e4b429606d02bc, }, // 0x14e4b429606d02bc3c604c0410e5fc01d6093a00bb3e2bc9395952af0b6a0dbd599a8782a1bea48a2aa4d8e1b1df7ca5
23  { 0x770553ef06aadf22, 0xc04c720ae01952ac, 0x9a55bee62edcdb27, 0x962f5650798fdf27, 0x28180e61b1f2cb8f, 0x0430df56ea4aba69, }, // 0x0430df56ea4aba6928180e61b1f2cb8f962f5650798fdf279a55bee62edcdb27c04c720ae01952ac770553ef06aadf22
24  },
25  g1p_x3 = {
26  { 0x7f0318e712da7559, 0x266ac7358912af30, 0xc4374a5888cfca69, 0x4f376827946368db, 0x61156a0c4d426519, 0x0bba0304fb3212a4, }, // 0x0bba0304fb3212a461156a0c4d4265194f376827946368dbc4374a5888cfca69266ac7358912af307f0318e712da7559
27  { 0xc25f47067af44e76, 0xcde09eab0276a0f4, 0xe9ec335b039fcf17, 0x727ce462858d3730, 0xdf4d86ed009f83fe, 0x0745d51f4d0912b3, }, // 0x0745d51f4d0912b3df4d86ed009f83fe727ce462858d3730e9ec335b039fcf17cde09eab0276a0f4c25f47067af44e76
28  { 0xceb902463027454d, 0xc9d68f804d8ec369, 0xc2ddd8e251d7339c, 0xd787b07101270da7, 0xfc3d86788b163753, 0x191fe9e914d73631, }, // 0x191fe9e914d73631fc3d86788b163753d787b07101270da7c2ddd8e251d7339cc9d68f804d8ec369ceb902463027454d
29  },
30  g1p_x24 = {
31  { 0xc8ff9a5471e72c92, 0x99683253e5aefa15, 0x5d6b135ab656eb43, 0x1b3776dc534fa4ab, 0xc2bfc4ab80c05017, 0x17b787b9910f9fa6, }, // 0x17b787b9910f9fa6c2bfc4ab80c050171b3776dc534fa4ab5d6b135ab656eb4399683253e5aefa15c8ff9a5471e72c92
32  { 0xf275ddf2d8723a25, 0xce36e492230ed9cd, 0xae724c9b9d46d006, 0x5d4cec21d5949cc3, 0x9ce9b30542ce5589, 0x05f77ff79a5b6f8a, }, // 0x05f77ff79a5b6f8a9ce9b30542ce55895d4cec21d5949cc3ae724c9b9d46d006ce36e492230ed9cdf275ddf2d8723a25
33  { 0x1eab6a2bf6bfeb17, 0xd88225cb44eaa0fb, 0x659281132d662bf8, 0xf0ac9c552dfd6f39, 0xb14437f70cc0f519, 0x0ae6513795046382, }, // 0x0ae6513795046382b14437f70cc0f519f0ac9c552dfd6f39659281132d662bf8d88225cb44eaa0fb1eab6a2bf6bfeb17
34  },
35  g1p_x25 = {
36  { 0x2e4f86255524abb3, 0xebb6095fb99f8e97, 0x3a6ab2001ab4f83c, 0x606df6ee661d3aa2, 0xf6b369b6a22b4047, 0x0b1416f427fc4c5f, }, // 0x0b1416f427fc4c5ff6b369b6a22b4047606df6ee661d3aa23a6ab2001ab4f83cebb6095fb99f8e972e4f86255524abb3
37  { 0xd0ee94cbc992ab24, 0x9fb5593bc61cd5bd, 0xc338a8acaef74389, 0x3a7da17eb290de91, 0xac616f60ea15f632, 0x0fc28d919f8ada25, }, // 0x0fc28d919f8ada25ac616f60ea15f6323a7da17eb290de91c338a8acaef743899fb5593bc61cd5bdd0ee94cbc992ab24
38  { 0x1bb5a6833e3677ae, 0xf3d50cd096cd2ceb, 0xa1d2c3cbc5527a6e, 0x60613c9426b3b9a1, 0xee0f3f71173f041c, 0x139ca4dd9f299816, }, // 0x139ca4dd9f299816ee0f3f71173f041c60613c9426b3b9a1a1d2c3cbc5527a6ef3d50cd096cd2ceb1bb5a6833e3677ae
39  };
40 
58 __global__ void G1TestKAT(testval_t *) {
59 
60  printf("=== RUN %s\n", __func__);
61 
62  bool pass = true;
63  size_t count = 0;
64 
65  g1p_t p, q, r, t, u, v;
66 
67  // 0=0
68 
69  if (pass) {
70 
71  g1p_inf(p); // p = 0
72 
73  if (g1p_neq(p, g1p_x0)) {
74  pass = false;
75 
76  g1p_print("0=0: FAIL: 0 != ", p);
77  }
78  ++count;
79  }
80 
81  // 0+0
82 
83  if (pass) {
84 
85  g1p_inf(p); // p = 0
86  g1p_inf(q); // q = 0
87  g1p_add(p, g1p_x0); // p += 0
88 
89  if (g1p_neq(p, q)) {
90  pass = false;
91 
92  g1p_print("0+0: FAIL: 0 != ", p);
93  }
94  ++count;
95  }
96 
97  // G+0 = 0+G = G
98 
99  if (pass) {
100 
101  g1p_inf(p); // p = 0
102  g1p_gen(q); // q = G
103  g1p_gen(r); // r = G
104 
105  g1p_add(p, q); // p += G
106  g1p_add(q, g1p_x0); // q += 0
107 
108  if (g1p_neq(p, r) || g1p_neq(q, r)) {
109  pass = false;
110 
111  g1p_print("0+G: = ", p);
112  g1p_print("G+0: = ", q);
113  g1p_print(" G: = ", r);
114  }
115  ++count;
116  }
117 
118  // G+G
119 
120  if (pass) {
121 
122  g1p_gen(p); // p = G
123  g1p_gen(q); // q = G
124  g1p_add(p, q); // p += q
125 
126  if (g1p_neq(p, g1p_x2)) {
127  pass = false;
128 
129  g1p_print("G+G: FAIL: 2G != ", p);
130  }
131  ++count;
132  }
133 
134  // 2G
135 
136  if (pass) {
137 
138  g1p_gen(p); // p = G
139  g1p_dbl(p); // p *= 2
140 
141  if (g1p_neq(p, g1p_x2)) {
142  pass = false;
143 
144  g1p_print("2G FAIL: 2G != ", p);
145  }
146  ++count;
147  }
148 
149  // G+2G
150 
151  if (pass) {
152 
153  g1p_gen(p); // p = G
154  g1p_add(p, g1p_x2); // p += 2G
155 
156  if (g1p_neq(p, g1p_x3)) {
157  pass = false;
158 
159  g1p_print("G+2G FAIL: 3G != ", p);
160  }
161  ++count;
162  }
163 
164  // 2*2*2*3G = 24G
165 
166  if (pass) {
167 
168  g1p_dbl(p); // 2G
169  g1p_dbl(p); // 2G
170  g1p_dbl(p); // 2G
171 
172  if (g1p_neq(p, g1p_x24)) {
173  pass = false;
174 
175  g1p_print("2*2*2*3G FAIL: 24G != ", p);
176  }
177  ++count;
178  }
179 
180  // 24G-2G+3G = 25G
181 
182  if (pass) {
183 
184  g1p_sub(p, g1p_x2); // 22G
185  g1p_add(p, g1p_x3); // 25G
186 
187  if (g1p_neq(p, g1p_x25)) {
188  pass = false;
189 
190  g1p_print("24G-2G+3G FAIL: 25G != ", p);
191  }
192  ++count;
193  }
194 
195  // 25 * G
196 
197  if (pass) {
198 
199  uint64_t x25[] = { 25, 0, 0, 0 };
200 
201  g1p_gen(p);
202  g1p_mul(p, x25); // 25G
203 
204  if (g1p_neq(p, g1p_x25)) {
205  pass = false;
206 
207  g1p_print("25G FAIL: 25G != ", p);
208  }
209  ++count;
210  }
211 
212  // 2G+G, 2G-G
213 
214  if (pass) {
215 
216  g1p_cpy(p, g1p_x2); // 2G
217  g1p_cpy(r, g1p_x2); // 2G
218 
219  g1p_gen(q); // 1G
220 
221  // Add-only reference
222  g1p_add(p, q);
223 
224  // Sub-only reference
225  g1p_sub(r, q);
226 
227  g1p_gen(p);
228  g1p_cpy(p, g1p_x2); // 2G
229 
230  // Add & subtract
231  g1p_addsub(p, q);
232 
233  if (g1p_neq(p, g1p_x3)) {
234  pass = false;
235  g1p_print("FAIL: 2G+G != ", p);
236  }
237 
238  g1p_gen(p);
239 
240  if (g1p_neq(q, p)) {
241  pass = false;
242 
243  g1p_print("FAIL: 2G-G != ", q);
244  }
245  ++count;
246  }
247 
248  if (pass) {
249 
250  g1p_gen(p); // 1G
251  g1p_gen(q); // 1G
252  g1p_gen(r); // 1G
253 
254  for (int i=0; pass && i<20000; i++) {
255  g1p_cpy(t, p);
256  g1p_cpy(u, q);
257  g1p_cpy(v, r);
258 
259  g1p_addsub(p, q);
260  g1p_addsub(p, q);
261  g1p_dbl(r);
262 
263  if (g1p_neq(p, q) || g1p_neq(q,r)) {
264  pass = false;
265 
266  printf("FAIL after %d ok:\n", i);
267  g1p_print("t = ", t);
268  g1p_print("u = ", u);
269  g1p_print("v = ", v);
270  g1p_print("p = ", p);
271  g1p_print("q = ", q);
272  g1p_print("r = ", r);
273  }
274  ++count;
275  }
276  }
277 
278  if ((blockIdx.x | blockIdx.y | blockIdx.z | threadIdx.x | threadIdx.y | threadIdx.z) == 0)
279  {
280  printf("%ld test%s\n", count, count == 1 ? "" : "s");
281 
282  PRINTPASS(pass);
283  }
284 }
285 
286 // vim: ts=4 et sw=4 si
g1p_addsub
__device__ void g1p_addsub(g1p_t &p, g1p_t &q)
Stores the sum and difference of p and q into p and q. Projective p and q, p,q ← p+q,...
Definition: g1p_addsub.cu:18
g1p_inf
__device__ __host__ void g1p_inf(g1p_t &p)
Set p to the point-at-infinity (0,1,0)
Definition: g1p.cu:93
g1p_sub
__device__ void g1p_sub(g1p_t &p, const g1p_t &q)
Point subtraction using projective coordinates. p ← p-q.
Definition: g1p_sub.cu:17
g1p_neq
__device__ bool g1p_neq(const g1p_t &p, const g1p_t &q)
Compares two projective points, returns true when not equal. This function compares if both parameter...
Definition: g1p_compare.cu:68
g1p_add
__device__ void g1p_add(g1p_t &p, const g1p_t &q)
Computes the sum of two points q into p, using projective coordinates. and stores in p.
Definition: g1p_add.cu:29
g1p_dbl
__device__ void g1p_dbl(g1p_t &p)
G1 point doubling, with write back: p=2*p.
Definition: g1p_dbl.cu:23
g1p_gen
__device__ __host__ void g1p_gen(g1p_t &p)
Sets p to the generator point G1 of bls12_381.
Definition: g1p.cu:106
g1p_mul
__device__ void g1p_mul(g1p_t &p, const fr_t &x)
p ← k·p Point multiplication by scalar, in projective coordinates. That result is stored back into p.
Definition: g1p_mul.cu:19
g1p_cpy
__device__ __host__ void g1p_cpy(g1p_t &p, const g1p_t &q)
Copy from q into p.
Definition: g1p.cu:67
g1p_print
__device__ __host__ void g1p_print(const char *s, const g1p_t &p)
Print a standard representation of p, preceded by the user-set string s.
Definition: g1p.cu:80
g1p_x2
__managed__ g1p_t g1p_x2
Definition: g1test_kat.cu:20
g1p_x25
__managed__ g1p_t g1p_x25
Definition: g1test_kat.cu:35
g1p_x3
__managed__ g1p_t g1p_x3
Definition: g1test_kat.cu:25
G1TestKAT
__global__ void G1TestKAT(testval_t *)
Test operation over G1 using KAT and self consistency:
Definition: g1test_kat.cu:58
g1p_x24
__managed__ g1p_t g1p_x24
Definition: g1test_kat.cu:30
g1p_x0
__managed__ g1p_t g1p_x0
Definition: g1test_kat.cu:11
g1p_t
G1 point in projective coordinates.
Definition: g1.cuh:27
testval_t
Definition: g1test.cuh:13
PRINTPASS
#define PRINTPASS(pass)
Definition: test.h:25