FK20 CUDA
frtest_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 "fr.cuh"
6 #include "frtest.cuh"
7 
15 __global__ void FrTestKAT(testval_t *) {
16 
17  printf("=== RUN %s\n", __func__);
18 
19  bool pass = true;
20  size_t count = 0;
21 
22  // fr_cpy
23 
24  if (pass) {
25  __const__ fr_t
26  q[] = {
27  { 0x9d419b78b22bbf1d, 0x69d3e1a350bcb984, 0x9382535ba5388d0f, 0x805e1591d721e13a },
28  { 0x32fa2230c45f3875, 0xdbd975eb0f0b10e3, 0xd69dbc9539ca9a98, 0x3c9fe7c9da36fc18 },
29  },
30  a[] = {
31  { 0x9d419b78b22bbf1d, 0x69d3e1a350bcb984, 0x9382535ba5388d0f, 0x805e1591d721e13a },
32  { 0x32fa2230c45f3875, 0xdbd975eb0f0b10e3, 0xd69dbc9539ca9a98, 0x3c9fe7c9da36fc18 },
33  };
34 
35  fr_t t;
36 
37  for (int i=0; pass && (i<2); i++) {
38  fr_cpy(t, q[i]);
39 
40  for (int j=0; j<4; j++)
41  if (t[j] != a[i][j])
42  pass = false;
43 
44  if (!pass)
45  printf("\n0x%016lx%016lx%016lx%016lx\n != 0x%016lx%016lx%016lx%016lx\nfr_cpy: FAIL\n",
46  t[3], t[2], t[1], t[0],
47  a[i][3], a[i][2], a[i][1], a[i][0]);
48 
49  ++count;
50  }
51  }
52 
53  // fr_reduce4
54 
55  if (pass) {
56  __const__ fr_t
57  q[] = {
58  { 0xFFFFFFFF00000000, 0x53BDA402FFFE5BFE, 0x3339D80809A1D805, 0x73EDA753299D7D48 }, // r-1
59  { 0xFFFFFFFF00000001, 0x53BDA402FFFE5BFE, 0x3339D80809A1D805, 0x73EDA753299D7D48 }, // r
60  { 0xFFFFFFFF00000002, 0x53BDA402FFFE5BFE, 0x3339D80809A1D805, 0x73EDA753299D7D48 }, // r+1
61  { 0xfffffffe00000001, 0xa77b4805fffcb7fd, 0x6673b0101343b00a, 0xe7db4ea6533afa90 }, // 2r-1
62  { 0xfffffffe00000002, 0xa77b4805fffcb7fd, 0x6673b0101343b00a, 0xe7db4ea6533afa90 }, // 2r
63  { 0xfffffffe00000003, 0xa77b4805fffcb7fd, 0x6673b0101343b00a, 0xe7db4ea6533afa90 }, // 2r+1
64  { 0xffffffffffffffff, 0xffffffffffffffff, 0xffffffffffffffff, 0xffffffffffffffff }, // 2^256-1
65  { 0xd2908bd8c20e727c, 0x894c03a208420535, 0x3d3db3045ad56524, 0x2a318cf9cca60638 },
66  { 0xe9ea6726afa1ac6a, 0x4a2e1c4163c05dca, 0x425e4a8e92fad8be, 0xdcec9f5991788284 },
67  { 0x7c9921ccff9fc38c, 0xdb84203d7e11de61, 0xa03712056f09f5b6, 0xba1104f592facb0d },
68  { 0x792a8035f135d8aa, 0xb5f66d21c0c92895, 0xa85a6dcdc58457d5, 0x755bb1f21e2fcba1 },
69  },
70  a[] = {
71  { 0xFFFFFFFF00000000, 0x53BDA402FFFE5BFE, 0x3339D80809A1D805, 0x73EDA753299D7D48 }, // r-1
72  { 0x0000000000000000, 0x0000000000000000, 0x0000000000000000, 0x0000000000000000 }, // r
73  { 0x0000000000000001, 0x0000000000000000, 0x0000000000000000, 0x0000000000000000 }, // r+1
74  { 0xFFFFFFFF00000000, 0x53BDA402FFFE5BFE, 0x3339D80809A1D805, 0x73EDA753299D7D48 }, // 2r-1
75  { 0x0000000000000000, 0x0000000000000000, 0x0000000000000000, 0x0000000000000000 }, // 2r
76  { 0x0000000000000001, 0x0000000000000000, 0x0000000000000000, 0x0000000000000000 }, // 2r+1
77  { 0x00000001fffffffd, 0x5884b7fa00034802, 0x998c4fefecbc4ff5, 0x1824b159acc5056f }, // 2^256-1
78  { 0xd2908bd8c20e727c, 0x894c03a208420535, 0x3d3db3045ad56524, 0x2a318cf9cca60638 },
79  { 0xe9ea6727afa1ac69, 0xf670783e63c201cb, 0x0f247286895900b8, 0x68fef80667db053c },
80  { 0x7c9921cdff9fc38b, 0x87c67c3a7e138262, 0x6cfd39fd65681db1, 0x46235da2695d4dc5 },
81  { 0x792a8036f135d8a9, 0x6238c91ec0cacc96, 0x752095c5bbe27fd0, 0x016e0a9ef4924e59 },
82  };
83 
84  fr_t t;
85 
86  for (int i=0; pass && (i<11); i++) {
87  fr_cpy(t, q[i]);
88 
89  fr_reduce4(t);
90 
91  for (int j=0; j<4; j++)
92  if (t[j] != a[i][j])
93  pass = false;
94 
95  if (!pass)
96  printf("\n0x%016lx%016lx%016lx%016lx\n != 0x%016lx%016lx%016lx%016lx\nfr_cpy: FAIL\n",
97  t[3], t[2], t[1], t[0],
98  a[i][3], a[i][2], a[i][1], a[i][0]);
99  ++count;
100  }
101  }
102 
103  // fr_eq, fr_neq
104 
105  if (pass) {
106  fr_t
107  t[] = {
108  { 0xFFFFFFFF00000000, 0x53BDA402FFFE5BFE, 0x3339D80809A1D805, 0x73EDA753299D7D48 }, // r-1
109  { 0xFFFFFFFF00000001, 0x53BDA402FFFE5BFE, 0x3339D80809A1D805, 0x73EDA753299D7D48 }, // r
110  { 0xFFFFFFFF00000002, 0x53BDA402FFFE5BFE, 0x3339D80809A1D805, 0x73EDA753299D7D48 }, // r+1
111  { 0xffffffffffffffff, 0xffffffffffffffff, 0xffffffffffffffff, 0xffffffffffffffff }, // 2^256-1
112  },
113  u[] = {
114  { 0xfffffffe00000001, 0xa77b4805fffcb7fd, 0x6673b0101343b00a, 0xe7db4ea6533afa90 }, // 2r-1
115  { 0xfffffffe00000002, 0xa77b4805fffcb7fd, 0x6673b0101343b00a, 0xe7db4ea6533afa90 }, // 2r
116  { 0xfffffffe00000003, 0xa77b4805fffcb7fd, 0x6673b0101343b00a, 0xe7db4ea6533afa90 }, // 2r+1
117  { 0x00000001fffffffd, 0x5884b7fa00034802, 0x998c4fefecbc4ff5, 0x1824b159acc5056f }, // 2^256-1
118  };
119 
120  for (int i=0; pass && i<4; i++) {
121  fr_t x;
122 
123  fr_cpy(x, t[i]);
124 
125  for (int j=0; pass && j<4; j++) {
126  fr_t y;
127 
128  fr_cpy(y, u[j]);
129 
130  int
131  eq = fr_eq (x, y),
132  neq = fr_neq(x, y);
133 
134  if (eq == neq) {
135  pass = false;
136 
137  printf("%d,%d: FAILED: inconsistent result, eq = %x, neq = %x\n", i, j, eq, neq);
138  }
139 
140  if ((i == j) && !eq) {
141  pass = false;
142 
143  printf("%d,%d: FAIL A: fr_eq claims inequality between these values:\n", i, j);
144 
145  printf("\t%016lX%016lX%016lX%016lX%016lX%016lX/\n",
146  t[i][5], t[i][4], t[i][3], t[i][2], t[i][1], t[i][0]);
147 
148  printf("\t%016lX%016lX%016lX%016lX%016lX%016lX\n",
149  x[5], x[4], x[3], x[2], x[1], x[0]);
150 
151  printf("\t%016lX%016lX%016lX%016lX%016lX%016lX/\n",
152  u[j][5], u[j][4], u[j][3], u[j][2], u[j][1], u[j][0]);
153 
154  printf("\t%016lX%016lX%016lX%016lX%016lX%016lX\n",
155  y[5], y[4], y[3], y[2], y[1], y[0]);
156 
157  printf("eq = %x, neq = %x\n", eq, neq);
158  }
159 
160  if ((i != j) && eq) {
161  pass = false;
162 
163  printf("%d,%d: FAIL B: fr_eq claims equality between these values:\n", i, j);
164 
165  printf("\t%016lX%016lX%016lX%016lX%016lX%016lX/\n",
166  t[i][5], t[i][4], t[i][3], t[i][2], t[i][1], t[i][0]);
167 
168  printf("\t%016lX%016lX%016lX%016lX%016lX%016lX\n",
169  x[5], x[4], x[3], x[2], x[1], x[0]);
170 
171  printf("\t%016lX%016lX%016lX%016lX%016lX%016lX/\n",
172  u[j][5], u[j][4], u[j][3], u[j][2], u[j][1], u[j][0]);
173 
174  printf("\t%016lX%016lX%016lX%016lX%016lX%016lX\n",
175  y[5], y[4], y[3], y[2], y[1], y[0]);
176 
177  printf("eq = %x, neq = %x\n", eq, neq);
178  }
179 
180  if ((i == j) && neq) {
181  pass = false;
182 
183  printf("%d,%d: FAIL C: fr_neq claims inequality between these values:\n", i, j);
184 
185  printf("\t%016lX%016lX%016lX%016lX%016lX%016lX/\n",
186  t[i][5], t[i][4], t[i][3], t[i][2], t[i][1], t[i][0]);
187 
188  printf("\t%016lX%016lX%016lX%016lX%016lX%016lX\n",
189  x[5], x[4], x[3], x[2], x[1], x[0]);
190 
191  printf("\t%016lX%016lX%016lX%016lX%016lX%016lX/\n",
192  u[j][5], u[j][4], u[j][3], u[j][2], u[j][1], u[j][0]);
193 
194  printf("\t%016lX%016lX%016lX%016lX%016lX%016lX\n",
195  y[5], y[4], y[3], y[2], y[1], y[0]);
196 
197  printf("eq = %x, neq = %x\n", eq, neq);
198  }
199 
200  if ((i != j) && !neq) {
201  pass = false;
202 
203  printf("%d,%d: FAIL D: fr_neq claims equality between these values:\n", i, j);
204 
205  printf("\t%016lX%016lX%016lX%016lX%016lX%016lX/\n",
206  t[i][5], t[i][4], t[i][3], t[i][2], t[i][1], t[i][0]);
207 
208  printf("\t%016lX%016lX%016lX%016lX%016lX%016lX\n",
209  x[5], x[4], x[3], x[2], x[1], x[0]);
210 
211  printf("\t%016lX%016lX%016lX%016lX%016lX%016lX/\n",
212  u[j][5], u[j][4], u[j][3], u[j][2], u[j][1], u[j][0]);
213 
214  printf("\t%016lX%016lX%016lX%016lX%016lX%016lX\n",
215  y[5], y[4], y[3], y[2], y[1], y[0]);
216 
217  printf("eq = %x, neq = %x\n", eq, neq);
218  }
219  ++count;
220  }
221  }
222  }
223 
224  // fr_neg
225 
226  if (pass) {
227  fr_t
228  q[] = {
229  { 0xfffffffe00000001, 0xa77b4805fffcb7fd, 0x6673b0101343b00a, 0xe7db4ea6533afa90 }, // 2p-1
230  { 0xfffffffe00000002, 0xa77b4805fffcb7fd, 0x6673b0101343b00a, 0xe7db4ea6533afa90 }, // 2p
231  { 0xfffffffe00000003, 0xa77b4805fffcb7fd, 0x6673b0101343b00a, 0xe7db4ea6533afa90 }, // 2p+1
232  { 0xffffffffffffffff, 0xffffffffffffffff, 0xffffffffffffffff, 0xffffffffffffffff }, // 2^256-1
233  { 0x8cc475b23946cd47, 0x59d85eba877626b1, 0x4a29ccf54ce7cf87, 0xcd2c630dbcf4adc7 },
234  { 0x602d7f007c75ea20, 0xa3548f2e30ddd9e5, 0xba434b1f4bf1f20a, 0xd022eb768da93eef },
235  { 0x85ac30859e8635ac, 0x27ebab61ee73762a, 0xd4343677f1ce5870, 0xa7b2394c1e305a80 },
236  { 0x5cf091ced52b5918, 0xbb40a07f7d360d4b, 0xbdc07041d2317755, 0x17d1f4993f9bfd51 },
237  },
238  a[] = {
239  { 0x0000000000000001, 0x0000000000000000, 0x0000000000000000, 0x0000000000000000 }, // r+1
240  { 0x0000000000000000, 0x0000000000000000, 0x0000000000000000, 0x0000000000000000 }, // r
241  { 0xffffffff00000000, 0x53bda402fffe5bfe, 0x3339d80809a1d805, 0x73eda753299d7d48 }, // r-1
242  { 0xfffffffc00000005, 0x4ef6900bfff96ffb, 0xcce7602026876015, 0xcfb69d4ca675f520 }, // 3p-2^256+1
243  { 0x733b8a4bc6b932bb, 0x4da2e94b7886914c, 0x1c49e31ac65be083, 0x1aaeeb9896464cc9 },
244  { 0x9fd280fd838a15e2, 0x0426b8d7cf1ede18, 0xac3064f0c751be00, 0x17b8632fc591bba0 },
245  { 0x7a53cf786179ca56, 0x7f8f9ca4118941d3, 0x923f79982175579a, 0x4029155a350aa00f },
246  { 0xa30f6e302ad4a6e9, 0x987d038382c84eb3, 0x757967c6377060af, 0x5c1bb2b9ea017ff6 },
247  };
248 
249  fr_t t;
250 
251  for (int i=0; pass && (i<8); i++) {
252  fr_cpy(t, q[i]);
253 
254  fr_neg(t);
255 
256  if (fr_neq(t, a[i]))
257  pass = false;
258 
259  if (!pass) {
260  printf("fr_neg: FAIL\n");
261 
262  printf("%d: 0x%016lx%016lx%016lx%016lx\n", i,
263  q[i][3], q[i][2], q[i][1], q[i][0]);
264 
265  fr_print("Expected ", a[i]);
266  fr_print("Received ", t);
267  }
268 
269  ++count;
270  }
271  }
272 
273  // fr_x2
274 
275  if (pass) {
276  fr_t
277  q[] = {
278  { 0xffffffff00000000, 0x53bda402fffe5bfe, 0x3339d80809a1d805, 0x73eda753299d7d48 }, // r-1
279  { 0xffffffff00000001, 0x53bda402fffe5bfe, 0x3339d80809a1d805, 0x73eda753299d7d48 }, // r
280  { 0xffffffff00000002, 0x53bda402fffe5bfe, 0x3339d80809a1d805, 0x73eda753299d7d48 }, // r+1
281  { 0xffffffffffffffff, 0xffffffffffffffff, 0xffffffffffffffff, 0xffffffffffffffff }, // 2^256-1
282  { 0xc88b1a71b810d7b4, 0x5cf286851064fbc2, 0x993b9498361cc386, 0xc4d51fa950015f9f },
283  { 0x60e55c7dd5623b5c, 0x33a1d44fb90ef709, 0xdc186fad1eaf837e, 0xfa211a6945f9fe12 },
284  { 0x2a05f2556d58ede0, 0xd5347f13e6e27df5, 0x05276cfe630c9247, 0x524f0ac0e3bb7b9d },
285  { 0xe2b5ffe82724e8f5, 0x47c1ec58ef0c159a, 0xdb1ccaa41fa78db2, 0x04c99214cc172e77 },
286  },
287  a[] = {
288  { 0xfffffffe00000000, 0xa77b4805fffcb7fd, 0x6673b0101343b00a, 0xe7db4ea6533afa90 }, // 2r-2
289  { 0xfffffffe00000002, 0xa77b4805fffcb7fd, 0x6673b0101343b00a, 0xe7db4ea6533afa90 }, // 2r
290  { 0xfffffffe00000004, 0xa77b4805fffcb7fd, 0x6673b0101343b00a, 0xe7db4ea6533afa90 }, // 2r+2
291  { 0x00000003fffffffa, 0xb1096ff400069004, 0x33189fdfd9789fea, 0x304962b3598a0adf }, // 2^257-2
292  { 0x911634e67021af65, 0xbeac210120cee388, 0x98c9a1184f53fefc, 0x2de14959232a4766 },
293  { 0xc1cab8ffaac476b4, 0x184d189372247e16, 0xeb497f3a16d7a6e7, 0x248b9785e57e0704 },
294  { 0x540be4abdab1dbbf, 0x56ab5a24cdc69feb, 0xd71501f4bc774c8a, 0x30b06e2e9dd979f1 },
295  { 0xc56bffd04e49d1ea, 0x8f83d8b1de182b35, 0xb63995483f4f1b64, 0x09932429982e5cef },
296  };
297 
298  fr_t t;
299 
300  for (int i=0; pass && (i<8); i++) {
301  fr_cpy(t, q[i]);
302 
303  fr_x2(t);
304 
305  if (fr_neq(t, a[i]))
306  pass = false;
307 
308  if (!pass) {
309  printf("fr_x2: FAIL\n");
310  printf("2 * 0x%016lx%016lx%016lx%016lx\n",
311  q[i][3], q[i][2], q[i][1], q[i][0]);
312 
313  fr_print("Expected ", a[i]);
314  fr_print("Received ", t);
315  }
316 
317  ++count;
318  }
319  }
320 
321  // fr_x3
322 
323  if (pass) {
324  fr_t
325  q[] = {
326  { 0xffffffff00000000, 0x53bda402fffe5bfe, 0x3339d80809a1d805, 0x73eda753299d7d48 }, // r-1
327  { 0xffffffff00000001, 0x53bda402fffe5bfe, 0x3339d80809a1d805, 0x73eda753299d7d48 }, // r
328  { 0xffffffff00000002, 0x53bda402fffe5bfe, 0x3339d80809a1d805, 0x73eda753299d7d48 }, // r+1
329  { 0xffffffffffffffff, 0xffffffffffffffff, 0xffffffffffffffff, 0xffffffffffffffff }, // 2^256-1
330  { 0x57dfda7aa31e3936, 0x09c6c5b91889a015, 0x78fc7321a936a1e5, 0xad831d7fec61991a },
331  { 0xb77c61179fa1d259, 0x677a048e37480e6f, 0x3e432db3a97d45bd, 0x1d7c65d774b61a0f },
332  { 0xc30ce006eb7c626f, 0xa9ec82098b214407, 0x19a27683d6e13e1c, 0xe2193f2301ead151 },
333  { 0x336eda9dd2484c45, 0xda56edce88b8077a, 0x5696673b95c09810, 0x636811703bb2ce69 },
334  },
335  a[] = {
336  { 0xfffffffdffffffff, 0xa77b4805fffcb7fd, 0x6673b0101343b00a, 0xe7db4ea6533afa90 }, // 2p-3
337  { 0xfffffffe00000002, 0xa77b4805fffcb7fd, 0x6673b0101343b00a, 0xe7db4ea6533afa90 }, // 2p
338  { 0xfffffffe00000005, 0xa77b4805fffcb7fd, 0x6673b0101343b00a, 0xe7db4ea6533afa90 }, // 2p+3
339  { 0x00000005fffffff7, 0x098e27ee0009d806, 0xcca4efcfc634efe0, 0x486e140d064f104e }, // 3*2^256-3
340  { 0x079f8f73e95aab9e, 0xce5dc11f49a37044, 0x9e0df944d51c8599, 0x38d2bb331eaed62e },
341  { 0x26752346dee5770b, 0x366e0daaa5d82b4f, 0xbac9891afc77d138, 0x587531865e224e2d },
342  { 0x4926a019c2752748, 0x5b11520da16c001c, 0x4cc62b63547a823b, 0x62a778c935ad018a },
343  { 0x9a4c8fdb76d8e4cd, 0xe78981659a2b5e70, 0x9d4f85a2adfe1827, 0x425ce5aa5fdd70ab },
344  };
345 
346  fr_t t;
347 
348  for (int i=0; pass && (i<8); i++) {
349  fr_cpy(t, q[i]);
350 
351  fr_x3(t);
352 
353  if (fr_neq(t, a[i]))
354  pass = false;
355 
356  if (!pass) {
357  printf("fr_x3: FAIL\n");
358  printf("3 * 0x%016lx%016lx%016lx%016lx\n",
359  q[i][3], q[i][2], q[i][1], q[i][0]);
360 
361  fr_print("Expected ", a[i]);
362  fr_print("Received ", t);
363  }
364 
365  ++count;
366  }
367  }
368 
369  // fr_add
370 
371  if (pass) {
372  fr_t
373  x[] = {
374  { 0xa9adcd7739bd74cf, 0xb53be4eb9f0aba5b, 0xf39db86dd74ba55e, 0x205eeb7c38dbbf9d },
375  { 0x1228a231f10f3969, 0x2f2fb5eebf0bbbe7, 0x0712c334b66a5db8, 0xebc3269a81507575 },
376  { 0xb600a76ec90fb9ed, 0x07e79b33160b873e, 0x5a9d273d378355e4, 0x44583bcadb79bbb9 },
377  { 0x504eb835fcb63c60, 0xbb06a2547d248706, 0xe9c922679104943d, 0xff4b832001a36d0d },
378  { 0xe31f46556875472c, 0xaed37a617a390c6a, 0x3566711770d0337e, 0xcab450c115f31533 },
379  { 0xccc5860e7b21e1c6, 0x58cadad8d1961158, 0x24fc2d78e8c1dd93, 0xc016a5c51c5a6199 },
380  { 0x1f5ceaaf4f7d6de4, 0x4c2caf2b9b11fa95, 0xbff1464b3d3d6939, 0x795aa4384f5ab137 },
381  { 0xf74a8f42e7f91926, 0x068815eb1cb349ca, 0xea0f7df125f638d4, 0x05dfc520280b525e },
382  },
383  y[] = {
384  { 0xc4fec1d957afa21a, 0x620cdaf477088c3e, 0x81c14702fd352d79, 0x579c2cab14ba5662 },
385  { 0x49031305fd08b422, 0x95247c6c292eafce, 0x66e5a460cc82bc8c, 0x8a6620ebfcebd51b },
386  { 0xd7a40e39d1285b4b, 0xf2ea148ea21f8495, 0x33ccff2766b4f4c9, 0x719912e5a5aba394 },
387  { 0xb0a9f4c1bd067f9e, 0xe57584bd59781e9a, 0x19ac2bbd2a9df826, 0x50a27f38bf3302ca },
388  { 0x6ce53032663eb897, 0x00b44b033aff2e55, 0xa1bb974e40d1e8df, 0x0cab9a0b906628af },
389  { 0x134aa1d8969bbf42, 0x8c36e257243dfcc3, 0x81cfbe199c2e992a, 0x797cd2e975ffcc1d },
390  { 0x7db61267bd4d87f2, 0x364422505bd94d9e, 0x3196afb585723831, 0x900e218d7c6fa297 },
391  { 0xa801fa7b481d5cca, 0x5e6809f0b1c9659f, 0x0087ad3b21681d47, 0xdf54afe9bc80f448 },
392  },
393  a[] = {
394  { 0x6eac8f51916d16e8, 0xc38b1bdd1614ea9b, 0x42252768cadefad2, 0x040d70d423f898b8 },
395  { 0x5b2bb53aee17ed88, 0xc91b4651e83f57b8, 0xd44adf7d66079234, 0x1a60518d0163d2b7 },
396  { 0x8da4b5a99a381537, 0xa7140bbeb82cafd5, 0x5b304e5c949672a8, 0x4203a75d5787e205 },
397  { 0x00f8acf9b9bcbbfc, 0xf900df0bd69feda3, 0x9d019e14a85edc59, 0x6812b3b26d9b7547 },
398  { 0x50047688ceb3ffc2, 0x5bca2161b539dec1, 0xa3e8305da8004458, 0x637243797cbbc09a },
399  { 0xe01027e911bda106, 0x3d867529f5d7561d, 0x40583b8271acc6b3, 0x51b82a083f1f3326 },
400  { 0x9d12fd190ccaf5d4, 0xdaf58975f6ee9035, 0x8b1445f0af6bf15f, 0x218d771f788f593e },
401  { 0x9f4c89bf301675ef, 0x11327bd8ce7e536b, 0xb75d53243dbc7e16, 0x7146cdb6baeec95e },
402  };
403 
404  fr_t t;
405 
406  for (int i=0; pass && (i<8); i++) {
407  fr_cpy(t, x[i]);
408  fr_add(t, y[i]);
409 
410  if (fr_neq(t, a[i])) {
411  pass = false;
412 
413  printf("fr_add: FAIL\n");
414  printf("0x%016lx%016lx%016lx%016lx + 0x%016lx%016lx%016lx%016lx\n",
415  x[i][3], x[i][2], x[i][1], x[i][0],
416  y[i][3], y[i][2], y[i][1], y[i][0]);
417 
418  fr_print("Expected ", a[i]);
419  fr_print("Received ", t);
420  }
421 
422  ++count;
423  }
424  }
425 
426  // fr_sub
427 
428  if (pass) {
429  fr_t
430  x[] = {
431  { 0x79c2dcd6c072a0eb, 0x5938348c63e5ee03, 0x2571278e41d8d960, 0x5cd48d8914135670 },
432  { 0xdd0e7887bd1fae7d, 0x496ce21047d6852e, 0x6257ac6e44173664, 0x2c7ec7596a88e143 },
433  { 0x4d1ae12876563e2f, 0x3544b4a08022241b, 0xdd5ac0fc26829b45, 0xde8c9d728a65a66b },
434  { 0xe42f83194305299d, 0x2bf8fc83ec477308, 0x15c5d82b5e0d969b, 0x4549a8ae17749941 },
435  { 0xcf43cbed7acb5065, 0xa62b1cef87f16d08, 0xce9aeb581a6f2561, 0x0aa57ffbe3f4eda8 },
436  { 0xd76412152770edf7, 0xf28e9ae385c7b25e, 0x3ce8377fe91310d9, 0x76fb671eeccdab6f },
437  { 0xe69ee87aff35023d, 0x1153b798996a5f4c, 0xde5d2875e048c920, 0x9932233eb0565fcd },
438  { 0xa9d696814e42f2d7, 0x803049704c5e877a, 0xc4873604dbd43251, 0xe5096ba7b364ef06 },
439  },
440  y[] = {
441  { 0x0525d4660db16d22, 0x5676c850dc0abe6a, 0x01d9e767cde7761d, 0x6adb7627e9f6aeef },
442  { 0x00b82535d67c3d30, 0x4726bb54e06ddf10, 0x4e6c6c189328a37c, 0x500f5f5b67ab1495 },
443  { 0x1a85df184987b062, 0x1f46f673e00873a1, 0x402323c56b47ff4d, 0xb56748bca526c301 },
444  { 0xf2adc10711825b80, 0xa015fddb814bb5cc, 0x1170a2b515955e6e, 0x69ad05e1728edb01 },
445  { 0x159ad40623ff267e, 0x9fee9451065ba006, 0xe7a0b5a0290501f7, 0x15225553b867d6ae },
446  { 0x2706fb06c7aa6448, 0x02a555b27cc700f8, 0x2c3994f637e68acb, 0x7c27cf5b5ea73d45 },
447  { 0x05aa619987706308, 0x9f3f454d588f659b, 0x17ed638bbc19695b, 0xccbf478a6c2b8bec },
448  { 0x1e827c98a11f5038, 0xc8ea30df2b736be2, 0x280a417b3dcd44d1, 0x85f8ca724eb03970 },
449  },
450  a[] = {
451  { 0x749d086fb2c133ca, 0x567f103e87d98b98, 0x56d1182e7d933b48, 0x65e6beb453ba24c9 },
452  { 0xdc565350e6a3714e, 0x5603cabe6767021d, 0x4725185dba906aed, 0x505d0f512c7b49f6 },
453  { 0x329502102cce8dcd, 0x15fdbe2ca019b07a, 0x9d379d36bb3a9bf8, 0x292554b5e53ee36a },
454  { 0xf181c2113182ce1e, 0xdfa0a2ab6afa193a, 0x378f0d7e521a1031, 0x4f8a4a1fce833b88 },
455  { 0xb9a8f7e656cc29e8, 0x59fa2ca181942901, 0x1a340dbffb0bfb6f, 0x6970d1fb552a9442 },
456  { 0xb05d170d5fc689b0, 0x43a6e93408ff0d65, 0x43e87a91bace5e14, 0x6ec13f16b7c3eb72 },
457  { 0xe0f486e077c49f36, 0xc5d2164e40d955b0, 0xf9a99cf22dd137c9, 0x406083076dc85129 },
458  { 0x8b5419e8ad23a29f, 0xb746189120eb1b98, 0x9c7cf4899e06ed7f, 0x5f10a13564b4b596 },
459  };
460 
461  fr_t t;
462 
463  for (int i=0; pass && (i<8); i++) {
464  fr_cpy(t, x[i]);
465 
466  fr_sub(t, y[i]);
467 
468  if (fr_neq(t, a[i]))
469  pass = false;
470 
471  if (!pass) {
472  printf("fr_sub: FAIL\n");
473  printf("0x%016lx%016lx%016lx%016lx%016lx%016lx - 0x%016lx%016lx%016lx%016lx%016lx%016lx\n",
474  x[i][5], x[i][4], x[i][3], x[i][2], x[i][1], x[i][0],
475  y[i][5], y[i][4], y[i][3], y[i][2], y[i][1], y[i][0]);
476 
477  fr_print("Expected ", a[i]);
478  fr_print("Received ", t);
479  }
480 
481  ++count;
482  }
483  }
484 
485  // fr_sqr
486 
487  if (pass) {
488  fr_t
489  q[] = {
490  { 0x8a28bc7a23b8ca27, 0x2d5a18614d5d87dc, 0x0b5170efac5b0908, 0x5d142f76332013e5 },
491  { 0x9e34a7d7bfd03f35, 0x3f8e1db0ff84a316, 0xc9041e4e1182ee8e, 0x24a6e8f76c1c85af },
492  { 0x023f831fe9692b89, 0x8e0c0ea2e6efeccc, 0x4f7ef565a60c3557, 0x25407371bb909e31 },
493  { 0x03f9bd4ad96916c6, 0xeeeb86b0e74a3539, 0x2877a95e76ce586d, 0x4d0458e170dbdd88 },
494  { 0x7cd618ac00921af4, 0x2d5ef5cbd3296ec9, 0xc91d64d567eeeb7a, 0x4d30d5d0b742a756 },
495  { 0x121e682e451a272c, 0x242abddbb631b4d1, 0x20e59d0cd3431333, 0xef33e5777473b021 },
496  { 0xee35cf214032be18, 0x81248cd6767199a1, 0xaabeba2bc6b6700c, 0x16c0e36fc3e8a70a },
497  { 0x2e0a1d0a32832d8e, 0xe7f1740cb372b6b8, 0xe92045e05b0b1ae1, 0x4f4034af157ec725 },
498  },
499  a[] = {
500  { 0xead982e107652210, 0x46dd627d4957c30e, 0xa85d9d7f02a8b745, 0x4b38e9a5e46488a3 },
501  { 0xc80b548070daf52c, 0x8eb33df3acf0542f, 0xb8b8b20de4b3d213, 0x034e5e23bf2e943f },
502  { 0x8266bc13d1addffd, 0x77b3de68740601b0, 0x65db94345d920468, 0x46d10123c4906c55 },
503  { 0x21a3dd420cff4523, 0x9d658b6044b8fb2a, 0x2e7e57b3eab308d1, 0x49a5d04eaea20b2e },
504  { 0x946fd1562247103c, 0x820fa121380274bc, 0x209dccc4a056dd09, 0x5d162117e34178a8 },
505  { 0x6c923642873be7d3, 0x39d5a34d22e476e3, 0xf9f58bacf3f1d0f3, 0x1f5926b571b05807 },
506  { 0xdc3eb6dacd33e8fa, 0x105b0f90c65e9c36, 0x2b3a600e428534ce, 0x04eca66fa60bd82c },
507  { 0x6c88a33e467d0436, 0x68de5bc32cb4b83d, 0x87e0d71ed13a8549, 0x0c6657a5ee1898ca },
508  };
509 
510  fr_t t;
511 
512  for (int i=0; pass && (i<8); i++) {
513  fr_cpy(t, q[i]);
514 
515  fr_sqr(t);
516 
517  if (fr_neq(t, a[i]))
518  pass = false;
519 
520  if (!pass) {
521  printf("fr_sqr: FAIL\n");
522  printf("0x%016lx%016lx%016lx%016lx^2\n",
523  q[i][3], q[i][2], q[i][1], q[i][0]);
524 
525  fr_print("Expected ", a[i]);
526  fr_print("Received ", t);
527  }
528 
529  ++count;
530  }
531  }
532 
533  // fr_mul
534 
535  if (pass) {
536  fr_t
537  x[] = {
538  { 0x3b44a96c59c690bb, 0x8697e87a732713a0, 0x945bea4e11648d71, 0x0c0aeeaacd9fdd46 },
539  { 0x3b68c9ecd7f475dc, 0x8fcbf94cc47a7712, 0x98262bc012006f6b, 0xcc59b69c6a11800e },
540  { 0x1c11a0b514009531, 0xbc79f3857c2d3e80, 0xb1b9ad274c9a1565, 0xabf34c7ef950991a },
541  { 0x53b0394c29c2cced, 0xbd81e570393678b6, 0xc887cec9b33c9221, 0x9eb4079889d23a2b },
542  { 0x161eb7e58a70f1be, 0x4d376cdcc50075ea, 0x421e12948915e8cb, 0xa94fe04dd3633195 },
543  { 0x7796b5a9ce8ccc50, 0x89a5032329875fc0, 0xb37708d4e7c0695a, 0x55ad168e639da67c },
544  { 0x5182e3055754826c, 0x244b84b1405dff5b, 0x7926637820bdf363, 0x0c4c38939953c602 },
545  { 0x57c3d05787a089e4, 0x05923cbbec717106, 0xebb314273a513815, 0x68504f715af9afe4 },
546  { 0x8775f85fefd3a3f7, 0x3d31cc4a7d359614, 0x0cbd12859051d3be, 0x18a286425c6a4f2a },
547  },
548  y[] = {
549  { 0xa5acba7322aca411, 0x53166475e534bae2, 0x3c01f18e78a68d32, 0x2b26e44aea71c03d },
550  { 0x390a4072d4249068, 0xed40e399d16d6bad, 0x6f45d1101c20d7c8, 0xf3892c2f5a40184b },
551  { 0x772b7b7c9b803667, 0x8fb29be985c830d0, 0x40a28d850e1b986b, 0x0378847e6b358f44 },
552  { 0x4170f9de04091161, 0x637f375b0fc38465, 0x3558e57591312120, 0xe007639449409917 },
553  { 0xf25e79d416c742cf, 0x5c62c18fb9f86579, 0x0de4ea6a4912155d, 0xf8a3e25080ca559d },
554  { 0xb8885f99f0bf3dee, 0x52ad47b58cc1fd03, 0xd645c5c990f5fcbd, 0xecfd9c6f82a56ce5 },
555  { 0xa8580fe78a036375, 0xb0bf9f7b46b905c8, 0x990a16f11f2fdf81, 0x5a0a27df527959df },
556  { 0x8bc3d4b107836984, 0xc274e53e00b19348, 0x86fcdbc0577d64e0, 0x534a17e34d281cd5 },
557  { 0x0000000000000020, 0x0000000000000000, 0x0000000000000000, 0x0000000000000000 },
558  },
559  a[] = {
560  { 0xde192398090105ca, 0xe59404bcd11fa709, 0xd92001ee4c393c0a, 0x139b5547381bff91 },
561  { 0x2de6763e47381e72, 0xf94e924974a5b1df, 0xaa3357e41e101621, 0x3461aff32809ad73 },
562  { 0x4b71a9acca4e8875, 0xcddb78d0cef1b76b, 0x498a2f1c6ab8c71d, 0x12be85e3b1f369e7 },
563  { 0x6fea58bd903735bc, 0x44fe8b10a81a2cc2, 0xa5ab3007b5a7b6bd, 0x4dc3b2d38ae169d9 },
564  { 0x3458bda5e25533f0, 0x8285e39ea2cdc36e, 0x145c2714ea99f555, 0x307c61d5ebb80b26 },
565  { 0x83c19aee28b8e10b, 0xe8d1a598a5aa6d41, 0x2d9d4e82499c6082, 0x2f6125d87fec738a },
566  { 0x1bc82ec8de6de6f6, 0x34994fa34fd514cc, 0xd106c6560453cd77, 0x24e2d06b5aafd115 },
567  { 0xd35fd5461360911d, 0xedf0e88e620985e2, 0x7623693cae99b940, 0x03787c0416bdc9c6 },
568  { 0xeebf0c03fa747eda, 0xafc7b13da6bc9a96, 0x64474081d06f67a7, 0x5cbedc589398f590 },
569  };
570 
571  fr_t t;
572 
573  for (int i=0; pass && (i<9); i++) {
574  fr_cpy(t, x[i]);
575 
576  fr_mul(t, y[i]);
577 
578  if (fr_neq(t, a[i]))
579  pass = false;
580 
581  if (!pass) {
582  printf("fr_mul: FAIL\n");
583  printf("0x%016lx%016lx%016lx%016lx * 0x%016lx%016lx%016lx%016lx\n",
584  x[i][3], x[i][2], x[i][1], x[i][0],
585  y[i][3], y[i][2], y[i][1], y[i][0]);
586 
587  fr_print("Expected ", a[i]);
588  fr_print("Received ", t);
589  }
590 
591  ++count;
592  }
593  }
594 
595  // fr_inv
596 
597  if (pass) {
598  fr_t
599  q[] = {
600  { 0x0000000000000000, 0x0000000000000000, 0x0000000000000000, 0x0000000000000000 },
601  { 0x0000000000000001, 0x0000000000000000, 0x0000000000000000, 0x0000000000000000 },
602  { 0x0000000000000002, 0x0000000000000000, 0x0000000000000000, 0x0000000000000000 },
603  { 0x0000000000000003, 0x0000000000000000, 0x0000000000000000, 0x0000000000000000 },
604  { 0xa127da24753d1ea9, 0xef5d7270e6940fb1, 0xc810db766a4a9105, 0x33344fab14e8ed3e },
605  { 0x56131b9360c13570, 0x060d28a6bf60b67e, 0x9ad516f2aa1dbd56, 0x519d20d7a485985d },
606  { 0x1a89fc6e8c94daed, 0x4db2fe4cec35a193, 0x46b68ee421699ad4, 0x664a986da9e91629 },
607  { 0x7e26ae14d79bcb20, 0xca9b6bd2f2c0d033, 0x3cdef4136ccab1d9, 0x5039948ba7d60a02 },
608  },
609  a[] = {
610  { 0x0000000000000000, 0x0000000000000000, 0x0000000000000000, 0x0000000000000000 },
611  { 0x0000000000000001, 0x0000000000000000, 0x0000000000000000, 0x0000000000000000 },
612  { 0x7fffffff80000001, 0xa9ded2017fff2dff, 0x199cec0404d0ec02, 0x39f6d3a994cebea4 },
613  { 0xaaaaaaaa00000001, 0xe27e6d5755543d54, 0xccd13ab0066be558, 0x4d491a377113a8da },
614  { 0x8cc00615238ce9a7, 0xe375dc0dbf0af565, 0xdba6b5f60714253a, 0x437bab96cfabdb5c },
615  { 0x87e1447096dd8c82, 0x4ac035465169b630, 0xce62b1c33527e264, 0x5dd785ebc525afa2 },
616  { 0x1db477d061fd382b, 0x4fed3127cdf422da, 0x4f85b52899197b2f, 0x41491bd8b882aa93 },
617  { 0x6abb0a53dcaacf19, 0x8f9169fb559d59c1, 0xfc2bb52133daa733, 0x6476eac140621485 },
618  };
619 
620  fr_t t;
621 
622  for (int i=0; pass && (i<8); i++) {
623  fr_cpy(t, q[i]);
624 
625  fr_inv(t);
626 
627  for (int j=0; j<4; j++)
628  if (t[j] != a[i][j])
629  pass = false;
630 
631  if (!pass) {
632  printf("fr_inv: FAIL\n");
633 
634  printf("1 / 0x%016lx%016lx%016lx%016lx\n",
635  q[i][3], q[i][2], q[i][1], q[i][0]);
636 
637  fr_print("Expected ", a[i]);
638  fr_print("Received ", t);
639  }
640 
641  ++count;
642  }
643  }
644 
645  printf("%ld tests\n", count);
646 
647  PRINTPASS(pass);
648 }
649 
650 // vim: ts=4 et sw=4 si
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_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_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_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_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_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_cpy
__device__ __host__ void fr_cpy(fr_t &z, const fr_t &x)
Copy from x into z.
Definition: fr_cpy.cu:14
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_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_inv
__device__ void fr_inv(fr_t &z)
Calculates the multiplicative inverse of z, and stores back.
Definition: fr_inv.cu:33
FrTestKAT
__global__ void FrTestKAT(testval_t *)
tests using KAT for: fr_cpy, fr_reduce4, fr_eq, fr_neq, fr_neg, fr_x2, fr_x3, fr_add,...
Definition: frtest_kat.cu:15
testval_t
Definition: g1test.cuh:13
PRINTPASS
#define PRINTPASS(pass)
Definition: test.h:25