FK20 CUDA
fp_inv.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 "fp.cuh"
6 
33 __device__ void fp_inv(fp_t &z, const fp_t &x) {
34 
35  // p-2 = 4002409555221667393417789825735904156556882819939007885332058136124031650490837864442687629129015664037894272559785
36  // = 5 * 5683 * 8572387 * 60565473370971138564533387 * 271298243114458311795492131869447332014235303245210008608904867493862992995991
37 
38  uint64_t x1[6], x3[6];
39 
40  // 5: 101
41 
42  fp_cpy(x1, x);
43  fp_sqr(z, x);
44  fp_sqr(z, z);
45  fp_mul(x1, z, x1);
46 
47  // 5683: 1011000110011
48 
49  fp_sqr(z, x1);
50  fp_mul(x3, z, x1);
51 
52  fp_sqr(z, z);
53  fp_sqr(z, z);
54  fp_mul(z, z, x3);
55  fp_sqr(z, z);
56  fp_sqr(z, z);
57  fp_sqr(z, z);
58  fp_sqr(z, z);
59  fp_sqr(z, z);
60  fp_mul(z, z, x3);
61  fp_sqr(z, z);
62  fp_sqr(z, z);
63  fp_sqr(z, z);
64  fp_sqr(z, z);
65  fp_mul(x1, z, x3);
66 
67  // 8572387: 100000101100110111100011
68  // 11 1 11 11 11 11
69  // 11
70 
71  fp_sqr(z, x1);
72  fp_mul(x3, z, x1);
73 
74  fp_sqr(z, z);
75  fp_sqr(z, z);
76  fp_sqr(z, z);
77  fp_sqr(z, z);
78  fp_sqr(z, z);
79  fp_mul(z, z, x1);
80  fp_sqr(z, z);
81  fp_sqr(z, z);
82  fp_sqr(z, z);
83  fp_mul(z, z, x3);
84  fp_sqr(z, z);
85  fp_sqr(z, z);
86  fp_sqr(z, z);
87  fp_sqr(z, z);
88  fp_mul(z, z, x3);
89  fp_sqr(z, z);
90  fp_sqr(z, z);
91  fp_sqr(z, z);
92  fp_mul(z, z, x3);
93  fp_sqr(z, z);
94  fp_sqr(z, z);
95  fp_mul(z, z, x3);
96  fp_sqr(z, z);
97  fp_sqr(z, z);
98  fp_sqr(z, z);
99  fp_sqr(z, z);
100  fp_sqr(z, z);
101  fp_mul(x1, z, x3);
102 
103  // 60565473370971138564533387: 11001000011001001111001110001111101100000100111001111101010010111100111010110010001011
104  // 11 1 11 1 11 11 11 1 11 1 11 11 1 1 1 1 11 11 1 11 1 1 11
105  // 11 1 11 1 11 11 1
106 
107  fp_sqr(z, x1);
108  fp_mul(z, z, x1);
109  fp_cpy(x3, z);
110 
111  fp_sqr(z, z);
112  fp_sqr(z, z);
113  fp_sqr(z, z);
114  fp_mul(z, z, x1);
115  fp_sqr(z, z);
116  fp_sqr(z, z);
117  fp_sqr(z, z);
118  fp_sqr(z, z);
119  fp_sqr(z, z);
120  fp_sqr(z, z);
121  fp_mul(z, z, x3);
122  fp_sqr(z, z);
123  fp_sqr(z, z);
124  fp_sqr(z, z);
125  fp_mul(z, z, x1);
126  fp_sqr(z, z);
127  fp_sqr(z, z);
128  fp_sqr(z, z);
129  fp_sqr(z, z);
130  fp_mul(z, z, x3);
131  fp_sqr(z, z);
132  fp_sqr(z, z);
133  fp_mul(z, z, x3);
134  fp_sqr(z, z);
135  fp_sqr(z, z);
136  fp_sqr(z, z);
137  fp_sqr(z, z);
138  fp_mul(z, z, x3);
139  fp_sqr(z, z);
140  fp_mul(z, z, x1);
141  fp_sqr(z, z);
142  fp_sqr(z, z);
143  fp_sqr(z, z);
144  fp_sqr(z, z);
145  fp_sqr(z, z);
146  fp_mul(z, z, x3);
147  fp_sqr(z, z);
148  fp_sqr(z, z);
149  fp_mul(z, z, x3);
150  fp_sqr(z, z);
151  fp_mul(z, z, x1);
152  fp_sqr(z, z);
153  fp_sqr(z, z);
154  fp_sqr(z, z);
155  fp_mul(z, z, x3);
156  fp_sqr(z, z);
157  fp_sqr(z, z);
158  fp_sqr(z, z);
159  fp_sqr(z, z);
160  fp_sqr(z, z);
161  fp_sqr(z, z);
162  fp_mul(z, z, x1);
163  fp_sqr(z, z);
164  fp_sqr(z, z);
165  fp_sqr(z, z);
166  fp_sqr(z, z);
167  fp_mul(z, z, x3);
168  fp_sqr(z, z);
169  fp_mul(z, z, x1);
170  fp_sqr(z, z);
171  fp_sqr(z, z);
172  fp_sqr(z, z);
173  fp_sqr(z, z);
174  fp_mul(z, z, x3);
175  fp_sqr(z, z);
176  fp_sqr(z, z);
177  fp_mul(z, z, x3);
178  fp_sqr(z, z);
179  fp_mul(z, z, x1);
180  fp_sqr(z, z);
181  fp_sqr(z, z);
182  fp_mul(z, z, x1);
183  fp_sqr(z, z);
184  fp_sqr(z, z);
185  fp_mul(z, z, x1);
186  fp_sqr(z, z);
187  fp_sqr(z, z);
188  fp_sqr(z, z);
189  fp_mul(z, z, x1);
190  fp_sqr(z, z);
191  fp_sqr(z, z);
192  fp_sqr(z, z);
193  fp_mul(z, z, x3);
194  fp_sqr(z, z);
195  fp_sqr(z, z);
196  fp_mul(z, z, x3);
197  fp_sqr(z, z);
198  fp_sqr(z, z);
199  fp_sqr(z, z);
200  fp_sqr(z, z);
201  fp_mul(z, z, x3);
202  fp_sqr(z, z);
203  fp_mul(z, z, x1);
204  fp_sqr(z, z);
205  fp_sqr(z, z);
206  fp_mul(z, z, x1);
207  fp_sqr(z, z);
208  fp_sqr(z, z);
209  fp_sqr(z, z);
210  fp_mul(z, z, x3);
211  fp_sqr(z, z);
212  fp_sqr(z, z);
213  fp_sqr(z, z);
214  fp_mul(z, z, x1);
215  fp_sqr(z, z);
216  fp_sqr(z, z);
217  fp_sqr(z, z);
218  fp_sqr(z, z);
219  fp_mul(z, z, x1);
220  fp_sqr(z, z);
221  fp_sqr(z, z);
222  fp_sqr(z, z);
223  fp_mul(x1, z, x3);
224 
225  // 271298243114458311795492131869447332014235303245210008608904867493862992995991
226 
227  fp_sqr(z, x1);
228  fp_mul(x3, z, x1);
229 
230  fp_sqr(z, z);
231  fp_sqr(z, z);
232  fp_mul(z, z, x1);
233  fp_sqr(z, z);
234  fp_sqr(z, z);
235  fp_mul(z, z, x1);
236  fp_sqr(z, z);
237  fp_sqr(z, z);
238  fp_sqr(z, z);
239  fp_mul(z, z, x3);
240  fp_sqr(z, z);
241  fp_sqr(z, z);
242  fp_mul(z, z, x3);
243  fp_sqr(z, z);
244  fp_mul(z, z, x1);
245  fp_sqr(z, z);
246  fp_sqr(z, z);
247  fp_sqr(z, z);
248  fp_sqr(z, z);
249  fp_mul(z, z, x3);
250  fp_sqr(z, z);
251  fp_sqr(z, z);
252  fp_mul(z, z, x1);
253  fp_sqr(z, z);
254  fp_sqr(z, z);
255  fp_mul(z, z, x1);
256  fp_sqr(z, z);
257  fp_sqr(z, z);
258  fp_sqr(z, z);
259  fp_mul(z, z, x3);
260  fp_sqr(z, z);
261  fp_mul(z, z, x1);
262  fp_sqr(z, z);
263  fp_sqr(z, z);
264  fp_mul(z, z, x1);
265  fp_sqr(z, z);
266  fp_sqr(z, z);
267  fp_sqr(z, z);
268  fp_sqr(z, z);
269  fp_mul(z, z, x3);
270  fp_sqr(z, z);
271  fp_sqr(z, z);
272  fp_sqr(z, z);
273  fp_sqr(z, z);
274  fp_sqr(z, z);
275  fp_sqr(z, z);
276  fp_mul(z, z, x1);
277  fp_sqr(z, z);
278  fp_sqr(z, z);
279  fp_sqr(z, z);
280  fp_sqr(z, z);
281  fp_mul(z, z, x3);
282  fp_sqr(z, z);
283  fp_mul(z, z, x1);
284  fp_sqr(z, z);
285  fp_sqr(z, z);
286  fp_sqr(z, z);
287  fp_mul(z, z, x3);
288  fp_sqr(z, z);
289  fp_sqr(z, z);
290  fp_mul(z, z, x3);
291  fp_sqr(z, z);
292  fp_sqr(z, z);
293  fp_sqr(z, z);
294  fp_sqr(z, z);
295  fp_mul(z, z, x3);
296  fp_sqr(z, z);
297  fp_sqr(z, z);
298  fp_sqr(z, z);
299  fp_mul(z, z, x3);
300  fp_sqr(z, z);
301  fp_sqr(z, z);
302  fp_sqr(z, z);
303  fp_mul(z, z, x1);
304  fp_sqr(z, z);
305  fp_sqr(z, z);
306  fp_sqr(z, z);
307  fp_mul(z, z, x1);
308  fp_sqr(z, z);
309  fp_sqr(z, z);
310  fp_sqr(z, z);
311  fp_sqr(z, z);
312  fp_mul(z, z, x3);
313  fp_sqr(z, z);
314  fp_mul(z, z, x1);
315  fp_sqr(z, z);
316  fp_sqr(z, z);
317  fp_sqr(z, z);
318  fp_sqr(z, z);
319  fp_mul(z, z, x1);
320  fp_sqr(z, z);
321  fp_sqr(z, z);
322  fp_sqr(z, z);
323  fp_sqr(z, z);
324  fp_sqr(z, z);
325  fp_sqr(z, z);
326  fp_sqr(z, z);
327  fp_mul(z, z, x3);
328  fp_sqr(z, z);
329  fp_sqr(z, z);
330  fp_sqr(z, z);
331  fp_sqr(z, z);
332  fp_sqr(z, z);
333  fp_sqr(z, z);
334  fp_sqr(z, z);
335  fp_mul(z, z, x3);
336  fp_sqr(z, z);
337  fp_sqr(z, z);
338  fp_mul(z, z, x3);
339  fp_sqr(z, z);
340  fp_sqr(z, z);
341  fp_mul(z, z, x1);
342  fp_sqr(z, z);
343  fp_sqr(z, z);
344  fp_sqr(z, z);
345  fp_sqr(z, z);
346  fp_mul(z, z, x3);
347  fp_sqr(z, z);
348  fp_sqr(z, z);
349  fp_mul(z, z, x3);
350  fp_sqr(z, z);
351  fp_sqr(z, z);
352  fp_sqr(z, z);
353  fp_sqr(z, z);
354  fp_mul(z, z, x1);
355  fp_sqr(z, z);
356  fp_sqr(z, z);
357  fp_mul(z, z, x1);
358  fp_sqr(z, z);
359  fp_sqr(z, z);
360  fp_sqr(z, z);
361  fp_mul(z, z, x3);
362  fp_sqr(z, z);
363  fp_sqr(z, z);
364  fp_sqr(z, z);
365  fp_sqr(z, z);
366  fp_mul(z, z, x3);
367  fp_sqr(z, z);
368  fp_sqr(z, z);
369  fp_sqr(z, z);
370  fp_sqr(z, z);
371  fp_mul(z, z, x1);
372  fp_sqr(z, z);
373  fp_sqr(z, z);
374  fp_sqr(z, z);
375  fp_sqr(z, z);
376  fp_sqr(z, z);
377  fp_mul(z, z, x3);
378  fp_sqr(z, z);
379  fp_sqr(z, z);
380  fp_mul(z, z, x3);
381  fp_sqr(z, z);
382  fp_sqr(z, z);
383  fp_mul(z, z, x3);
384  fp_sqr(z, z);
385  fp_sqr(z, z);
386  fp_sqr(z, z);
387  fp_sqr(z, z);
388  fp_sqr(z, z);
389  fp_sqr(z, z);
390  fp_sqr(z, z);
391  fp_sqr(z, z);
392  fp_mul(z, z, x1);
393  fp_sqr(z, z);
394  fp_sqr(z, z);
395  fp_sqr(z, z);
396  fp_sqr(z, z);
397  fp_sqr(z, z);
398  fp_sqr(z, z);
399  fp_mul(z, z, x3);
400  fp_sqr(z, z);
401  fp_mul(z, z, x1);
402  fp_sqr(z, z);
403  fp_sqr(z, z);
404  fp_mul(z, z, x1);
405  fp_sqr(z, z);
406  fp_sqr(z, z);
407  fp_sqr(z, z);
408  fp_mul(z, z, x3);
409  fp_sqr(z, z);
410  fp_mul(z, z, x1);
411  fp_sqr(z, z);
412  fp_sqr(z, z);
413  fp_sqr(z, z);
414  fp_sqr(z, z);
415  fp_mul(z, z, x1);
416  fp_sqr(z, z);
417  fp_sqr(z, z);
418  fp_sqr(z, z);
419  fp_sqr(z, z);
420  fp_sqr(z, z);
421  fp_sqr(z, z);
422  fp_sqr(z, z);
423  fp_sqr(z, z);
424  fp_sqr(z, z);
425  fp_mul(z, z, x3);
426  fp_sqr(z, z);
427  fp_mul(z, z, x1);
428  fp_sqr(z, z);
429  fp_sqr(z, z);
430  fp_mul(z, z, x1);
431  fp_sqr(z, z);
432  fp_sqr(z, z);
433  fp_mul(z, z, x1);
434  fp_sqr(z, z);
435  fp_sqr(z, z);
436  fp_sqr(z, z);
437  fp_sqr(z, z);
438  fp_sqr(z, z);
439  fp_sqr(z, z);
440  fp_mul(z, z, x1);
441  fp_sqr(z, z);
442  fp_sqr(z, z);
443  fp_sqr(z, z);
444  fp_mul(z, z, x3);
445  fp_sqr(z, z);
446  fp_sqr(z, z);
447  fp_sqr(z, z);
448  fp_mul(z, z, x1);
449  fp_sqr(z, z);
450  fp_sqr(z, z);
451  fp_sqr(z, z);
452  fp_sqr(z, z);
453  fp_sqr(z, z);
454  fp_mul(z, z, x3);
455  fp_sqr(z, z);
456  fp_sqr(z, z);
457  fp_mul(z, z, x3);
458  fp_sqr(z, z);
459  fp_sqr(z, z);
460  fp_mul(z, z, x3);
461  fp_sqr(z, z);
462  fp_sqr(z, z);
463  fp_sqr(z, z);
464  fp_mul(z, z, x3);
465  fp_sqr(z, z);
466  fp_sqr(z, z);
467  fp_mul(z, z, x3);
468  fp_sqr(z, z);
469  fp_sqr(z, z);
470  fp_mul(z, z, x3);
471  fp_sqr(z, z);
472  fp_sqr(z, z);
473  fp_mul(z, z, x3);
474  fp_sqr(z, z);
475  fp_sqr(z, z);
476  fp_mul(z, z, x1);
477  fp_sqr(z, z);
478  fp_sqr(z, z);
479  fp_sqr(z, z);
480  fp_mul(z, z, x1);
481  fp_sqr(z, z);
482  fp_sqr(z, z);
483  fp_sqr(z, z);
484  fp_mul(z, z, x3);
485  fp_sqr(z, z);
486  fp_sqr(z, z);
487  fp_sqr(z, z);
488  fp_sqr(z, z);
489  fp_mul(z, z, x3);
490  fp_sqr(z, z);
491  fp_sqr(z, z);
492  fp_sqr(z, z);
493  fp_mul(z, z, x3);
494  fp_sqr(z, z);
495  fp_sqr(z, z);
496  fp_sqr(z, z);
497  fp_mul(z, z, x1);
498  fp_sqr(z, z);
499  fp_sqr(z, z);
500  fp_mul(z, z, x1);
501  fp_sqr(z, z);
502  fp_sqr(z, z);
503  fp_sqr(z, z);
504  fp_sqr(z, z);
505  fp_mul(z, z, x3);
506  fp_sqr(z, z);
507  fp_sqr(z, z);
508  fp_mul(z, z, x1);
509  fp_sqr(z, z);
510  fp_sqr(z, z);
511  fp_sqr(z, z);
512  fp_mul(z, z, x3);
513  fp_sqr(z, z);
514  fp_sqr(z, z);
515  fp_mul(z, z, x3);
516  fp_sqr(z, z);
517  fp_mul(z, z, x1);
518  fp_sqr(z, z);
519  fp_sqr(z, z);
520  fp_sqr(z, z);
521  fp_sqr(z, z);
522  fp_sqr(z, z);
523  fp_sqr(z, z);
524  fp_mul(z, z, x1);
525  fp_sqr(z, z);
526  fp_sqr(z, z);
527  fp_sqr(z, z);
528  fp_mul(z, z, x3);
529  fp_sqr(z, z);
530  fp_mul(z, z, x1);
531  fp_sqr(z, z);
532  fp_sqr(z, z);
533  fp_mul(z, z, x1);
534  fp_sqr(z, z);
535  fp_sqr(z, z);
536  fp_sqr(z, z);
537  fp_sqr(z, z);
538  fp_mul(z, z, x3);
539  fp_sqr(z, z);
540  fp_sqr(z, z);
541  fp_mul(z, z, x1);
542  fp_sqr(z, z);
543  fp_sqr(z, z);
544  fp_sqr(z, z);
545  fp_sqr(z, z);
546  fp_mul(z, z, x3);
547  fp_sqr(z, z);
548  fp_mul(z, z, x1);
549  fp_sqr(z, z);
550  fp_sqr(z, z);
551  fp_sqr(z, z);
552  fp_mul(z, z, x3);
553  fp_sqr(z, z);
554  fp_sqr(z, z);
555  fp_sqr(z, z);
556  fp_sqr(z, z);
557  fp_mul(z, z, x3);
558  fp_sqr(z, z);
559  fp_sqr(z, z);
560  fp_mul(z, z, x1);
561  fp_sqr(z, z);
562  fp_sqr(z, z);
563  fp_sqr(z, z);
564  fp_mul(z, z, x1);
565  fp_sqr(z, z);
566  fp_sqr(z, z);
567  fp_sqr(z, z);
568  fp_mul(z, z, x3);
569  fp_sqr(z, z);
570  fp_mul(z, z, x1);
571 }
572 
573 // vim: ts=4 et sw=4 si
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_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_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_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