FK20 CUDA
fr_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 "fr.cuh"
6 #include <stdio.h>
7 
33 __device__ void fr_inv(fr_t &z) {
34 
35  // r-2 = 52435875175126190479447740508185965837690552500527637822603658699938581184511
36  // = 111001111101101101001110101001100101001100111010111\
37  110101001000001100110011100111011000000010000000100\
38  110100001110110000000010101010011101111011010010000\
39  000010111111111111111001011011111111101111111111111\
40  111111111111111111011111111111111111111111111111111
41 
42  uint64_t x1[4], x3[4], x5[4], x7[4];
43 
44  fr_cpy(x1, z); // 1
45  fr_sqr(z); // 10
46 
47  fr_cpy(x3, x1); // 1
48  fr_mul(x3, z); // 11
49 
50  fr_cpy(x5, x3); // 11
51  fr_mul(x5, z); // 101
52 
53  fr_cpy(x7, x5); // 101
54  fr_mul(x7, z); // 111
55 
56  fr_cpy(z, x7); // 111
57 
58  fr_sqr(z);
59  fr_sqr(z);
60  fr_sqr(z);
61  fr_sqr(z);
62  fr_sqr(z);
63  fr_mul(z, x5); // 11100101
64 
65  fr_sqr(z);
66  fr_mul(z, x5); // 111001111
67 
68  fr_sqr(z);
69  fr_sqr(z);
70  fr_sqr(z);
71  fr_mul(z, x5); // 111001111101
72 
73  fr_sqr(z);
74  fr_sqr(z);
75  fr_sqr(z);
76  fr_mul(z, x5); // 111001111101101
77 
78  fr_sqr(z);
79  fr_sqr(z);
80  fr_sqr(z);
81  fr_mul(z, x5); // 111001111101101101
82 
83  fr_sqr(z);
84  fr_sqr(z);
85  fr_sqr(z);
86  fr_sqr(z);
87  fr_sqr(z);
88  fr_mul(z, x7); // 11100111110110110100111
89 
90  fr_sqr(z);
91  fr_sqr(z);
92  fr_sqr(z);
93  fr_sqr(z);
94  fr_mul(z, x5); // 111001111101101101001110101
95 
96  fr_sqr(z);
97  fr_sqr(z);
98  fr_sqr(z);
99  fr_sqr(z);
100  fr_mul(z, x3); // 1110011111011011010011101010011
101 
102  fr_sqr(z);
103  fr_sqr(z);
104  fr_sqr(z);
105  fr_sqr(z);
106  fr_sqr(z);
107  fr_mul(z, x5); // 111001111101101101001110101001100101
108 
109  fr_sqr(z);
110  fr_sqr(z);
111  fr_sqr(z);
112  fr_sqr(z);
113  fr_mul(z, x3); // 1110011111011011010011101010011001010011
114 
115  fr_sqr(z);
116  fr_sqr(z);
117  fr_sqr(z);
118  fr_sqr(z);
119  fr_sqr(z);
120  fr_mul(z, x7); // 111001111101101101001110101001100101001100111
121 
122  fr_sqr(z);
123  fr_sqr(z);
124  fr_sqr(z);
125  fr_sqr(z);
126  fr_mul(z, x5); // 1110011111011011010011101010011001010011001110101
127 
128  fr_sqr(z);
129  fr_sqr(z);
130  fr_sqr(z);
131  fr_mul(z, x7); // 1110011111011011010011101010011001010011001110101111
132 
133  fr_sqr(z);
134  fr_sqr(z);
135  fr_sqr(z);
136  fr_mul(z, x5); // 1110011111011011010011101010011001010011001110101111101
137 
138  fr_sqr(z);
139  fr_sqr(z);
140  fr_mul(z, x1); // 111001111101101101001110101001100101001100111010111110101
141 
142  fr_sqr(z);
143  fr_sqr(z);
144  fr_sqr(z);
145  fr_mul(z, x1); // 111001111101101101001110101001100101001100111010111110101001
146 
147  fr_sqr(z);
148  fr_sqr(z);
149  fr_sqr(z);
150  fr_sqr(z);
151  fr_sqr(z);
152  fr_sqr(z);
153  fr_sqr(z);
154  fr_mul(z, x3); // 1110011111011011010011101010011001010011001110101111101010010000011
155 
156  fr_sqr(z);
157  fr_sqr(z);
158  fr_sqr(z);
159  fr_sqr(z);
160  fr_mul(z, x3); // 11100111110110110100111010100110010100110011101011111010100100000110011
161 
162  fr_sqr(z);
163  fr_sqr(z);
164  fr_sqr(z);
165  fr_sqr(z);
166  fr_sqr(z);
167  fr_mul(z, x7); // 1110011111011011010011101010011001010011001110101111101010010000011001100111
168 
169  fr_sqr(z);
170  fr_sqr(z);
171  fr_sqr(z);
172  fr_sqr(z);
173  fr_sqr(z);
174  fr_mul(z, x7); // 111001111101101101001110101001100101001100111010111110101001000001100110011100111
175 
176  fr_sqr(z);
177  fr_sqr(z);
178  fr_sqr(z);
179  fr_mul(z, x3); // 111001111101101101001110101001100101001100111010111110101001000001100110011100111011
180 
181  fr_sqr(z);
182  fr_sqr(z);
183  fr_sqr(z);
184  fr_sqr(z);
185  fr_sqr(z);
186  fr_sqr(z);
187  fr_sqr(z);
188  fr_sqr(z);
189  fr_mul(z, x1); // 11100111110110110100111010100110010100110011101011111010100100000110011001110011101100000001
190 
191  fr_sqr(z);
192  fr_sqr(z);
193  fr_sqr(z);
194  fr_sqr(z);
195  fr_sqr(z);
196  fr_sqr(z);
197  fr_sqr(z);
198  fr_sqr(z);
199  fr_mul(z, x1); // 1110011111011011010011101010011001010011001110101111101010010000011001100111001110110000000100000001
200 
201  fr_sqr(z);
202  fr_sqr(z);
203  fr_sqr(z);
204  fr_sqr(z);
205  fr_mul(z, x3); // 11100111110110110100111010100110010100110011101011111010100100000110011001110011101100000001000000010011
206 
207  fr_sqr(z);
208  fr_sqr(z);
209  fr_mul(z, x1); // 1110011111011011010011101010011001010011001110101111101010010000011001100111001110110000000100000001001101
210 
211  fr_sqr(z);
212  fr_sqr(z);
213  fr_sqr(z);
214  fr_sqr(z);
215  fr_sqr(z);
216  fr_sqr(z);
217  fr_sqr(z);
218  fr_mul(z, x7); // 11100111110110110100111010100110010100110011101011111010100100000110011001110011101100000001000000010011010000111
219 
220  fr_sqr(z);
221  fr_sqr(z);
222  fr_sqr(z);
223  fr_mul(z, x3); // 11100111110110110100111010100110010100110011101011111010100100000110011001110011101100000001000000010011010000111011
224 
225  fr_sqr(z);
226  fr_sqr(z);
227  fr_sqr(z);
228  fr_sqr(z);
229  fr_sqr(z);
230  fr_sqr(z);
231  fr_sqr(z);
232  fr_sqr(z);
233  fr_sqr(z);
234  fr_sqr(z);
235  fr_sqr(z);
236  fr_mul(z, x5); // 1110011111011011010011101010011001010011001110101111101010010000011001100111001110110000000100000001001101000011101100000000101
237 
238  fr_sqr(z);
239  fr_sqr(z);
240  fr_sqr(z);
241  fr_sqr(z);
242  fr_mul(z, x5); // 11100111110110110100111010100110010100110011101011111010100100000110011001110011101100000001000000010011010000111011000000001010101
243 
244  fr_sqr(z);
245  fr_sqr(z);
246  fr_sqr(z);
247  fr_sqr(z);
248  fr_sqr(z);
249  fr_mul(z, x7); // 1110011111011011010011101010011001010011001110101111101010010000011001100111001110110000000100000001001101000011101100000000101010100111
250 
251  fr_sqr(z);
252  fr_sqr(z);
253  fr_sqr(z);
254  fr_sqr(z);
255  fr_mul(z, x7); // 11100111110110110100111010100110010100110011101011111010100100000110011001110011101100000001000000010011010000111011000000001010101001110111
256 
257  fr_sqr(z);
258  fr_sqr(z);
259  fr_sqr(z);
260  fr_mul(z, x5); // 11100111110110110100111010100110010100110011101011111010100100000110011001110011101100000001000000010011010000111011000000001010101001110111101
261 
262  fr_sqr(z);
263  fr_sqr(z);
264  fr_sqr(z);
265  fr_mul(z, x5); // 11100111110110110100111010100110010100110011101011111010100100000110011001110011101100000001000000010011010000111011000000001010101001110111101101
266 
267  fr_sqr(z);
268  fr_sqr(z);
269  fr_sqr(z);
270  fr_mul(z, x1); // 11100111110110110100111010100110010100110011101011111010100100000110011001110011101100000001000000010011010000111011000000001010101001110111101101001
271 
272  fr_sqr(z);
273  fr_sqr(z);
274  fr_sqr(z);
275  fr_sqr(z);
276  fr_sqr(z);
277  fr_sqr(z);
278  fr_sqr(z);
279  fr_sqr(z);
280  fr_sqr(z);
281  fr_sqr(z);
282  fr_sqr(z);
283  fr_mul(z, x5); // 1110011111011011010011101010011001010011001110101111101010010000011001100111001110110000000100000001001101000011101100000000101010100111011110110100100000000101
284 
285  // replace x7 with x127
286 
287  fr_sqr(x7);
288  fr_sqr(x7);
289  fr_mul(x7, x3); // 11111
290  fr_sqr(x7);
291  fr_sqr(x7);
292  fr_mul(x7, x3); // 1111111
293 
294  fr_sqr(z);
295  fr_sqr(z);
296  fr_sqr(z);
297  fr_sqr(z);
298  fr_sqr(z);
299  fr_sqr(z);
300  fr_sqr(z);
301  fr_mul(z, x7); // 11100111110110110100111010100110010100110011101011111010100100000110011001110011101100000001000000010011010000111011000000001010101001110111101101001000000001011111111
302 
303  fr_sqr(z);
304  fr_sqr(z);
305  fr_sqr(z);
306  fr_sqr(z);
307  fr_sqr(z);
308  fr_sqr(z);
309  fr_sqr(z);
310  fr_mul(z, x7); // 111001111101101101001110101001100101001100111010111110101001000001100110011100111011000000010000000100110100001110110000000010101010011101111011010010000000010111111111111111
311 
312  fr_sqr(z);
313  fr_sqr(z);
314  fr_sqr(z);
315  fr_sqr(z);
316  fr_sqr(z);
317  fr_mul(z, x5); // 11100111110110110100111010100110010100110011101011111010100100000110011001110011101100000001000000010011010000111011000000001010101001110111101101001000000001011111111111111100101
318 
319  fr_sqr(z);
320  fr_sqr(z);
321  fr_sqr(z);
322  fr_mul(z, x5); // 11100111110110110100111010100110010100110011101011111010100100000110011001110011101100000001000000010011010000111011000000001010101001110111101101001000000001011111111111111100101101
323 
324  fr_sqr(z);
325  fr_sqr(z);
326  fr_sqr(z);
327  fr_sqr(z);
328  fr_sqr(z);
329  fr_sqr(z);
330  fr_sqr(z);
331  fr_mul(z, x7); // 111001111101101101001110101001100101001100111010111110101001000001100110011100111011000000010000000100110100001110110000000010101010011101111011010010000000010111111111111111001011011111111
332 
333  fr_sqr(z);
334  fr_sqr(z);
335  fr_sqr(z);
336  fr_mul(z, x5); // 111001111101101101001110101001100101001100111010111110101001000001100110011100111011000000010000000100110100001110110000000010101010011101111011010010000000010111111111111111001011011111111101
337 
338  fr_sqr(z);
339  fr_sqr(z);
340  fr_sqr(z);
341  fr_sqr(z);
342  fr_sqr(z);
343  fr_sqr(z);
344  fr_sqr(z);
345  fr_mul(z, x7); // 1110011111011011010011101010011001010011001110101111101010010000011001100111001110110000000100000001001101000011101100000000101010100111011110110100100000000101111111111111110010110111111111011111111
346 
347  fr_sqr(z);
348  fr_sqr(z);
349  fr_sqr(z);
350  fr_sqr(z);
351  fr_sqr(z);
352  fr_sqr(z);
353  fr_sqr(z);
354  fr_mul(z, x7); // 11100111110110110100111010100110010100110011101011111010100100000110011001110011101100000001000000010011010000111011000000001010101001110111101101001000000001011111111111111100101101111111110111111111111111
355 
356  // replace x127 with x255
357 
358  fr_sqr(x7);
359  fr_mul(x7, x1); // 11111111
360 
361  fr_sqr(z);
362  fr_sqr(z);
363  fr_sqr(z);
364  fr_sqr(z);
365  fr_sqr(z);
366  fr_sqr(z);
367  fr_sqr(z);
368  fr_sqr(z);
369  fr_mul(z, x7); // 1110011111011011010011101010011001010011001110101111101010010000011001100111001110110000000100000001001101000011101100000000101010100111011110110100100000000101111111111111110010110111111111011111111111111111111111
370 
371  fr_sqr(z);
372  fr_sqr(z);
373  fr_sqr(z);
374  fr_sqr(z);
375  fr_sqr(z);
376  fr_sqr(z);
377  fr_sqr(z);
378  fr_sqr(z);
379  fr_mul(z, x7); // 111001111101101101001110101001100101001100111010111110101001000001100110011100111011000000010000000100110100001110110000000010101010011101111011010010000000010111111111111111001011011111111101111111111111111111111111111111
380 
381  fr_sqr(z);
382  fr_sqr(z);
383  fr_sqr(z);
384  fr_sqr(z);
385  fr_sqr(z);
386  fr_sqr(z);
387  fr_sqr(z);
388  fr_sqr(z);
389  fr_sqr(z);
390  fr_mul(z, x7); // 111001111101101101001110101001100101001100111010111110101001000001100110011100111011000000010000000100110100001110110000000010101010011101111011010010000000010111111111111111001011011111111101111111111111111111111111111111011111111
391 
392  fr_sqr(z);
393  fr_sqr(z);
394  fr_sqr(z);
395  fr_sqr(z);
396  fr_sqr(z);
397  fr_sqr(z);
398  fr_sqr(z);
399  fr_sqr(z);
400  fr_mul(z, x7); // 11100111110110110100111010100110010100110011101011111010100100000110011001110011101100000001000000010011010000111011000000001010101001110111101101001000000001011111111111111100101101111111110111111111111111111111111111111101111111111111111
401 
402  fr_sqr(z);
403  fr_sqr(z);
404  fr_sqr(z);
405  fr_sqr(z);
406  fr_sqr(z);
407  fr_sqr(z);
408  fr_sqr(z);
409  fr_sqr(z);
410  fr_mul(z, x7); // 1110011111011011010011101010011001010011001110101111101010010000011001100111001110110000000100000001001101000011101100000000101010100111011110110100100000000101111111111111110010110111111111011111111111111111111111111111110111111111111111111111111
411 
412  fr_sqr(z);
413  fr_sqr(z);
414  fr_sqr(z);
415  fr_sqr(z);
416  fr_sqr(z);
417  fr_sqr(z);
418  fr_sqr(z);
419  fr_sqr(z);
420  fr_mul(z, x7); // 111001111101101101001110101001100101001100111010111110101001000001100110011100111011000000010000000100110100001110110000000010101010011101111011010010000000010111111111111111001011011111111101111111111111111111111111111111011111111111111111111111111111111
421 }
422 
423 // vim: ts=4 et sw=4 si
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_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_cpy
__device__ __host__ void fr_cpy(fr_t &z, const fr_t &x)
Copy from x into z.
Definition: fr_cpy.cu:14
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