FK20 CUDA
frtest_mulconst.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 
8 #define ITER 80
9 
21 __global__ void FrTestMulConst(testval_t *testval) {
22 
23  printf("=== RUN %s\n", __func__);
24 
25  bool pass = true;
26  size_t count = 0;
27 
28  // 2*4 == 8
29 
30  for (int i=0; pass && i<TESTVALS; i++) {
31  fr_t x2x4, x8;
32 
33  fr_cpy(x2x4, testval[i]);
34  fr_cpy(x8, testval[i]);
35 
36  for (int j=0; pass && j<ITER; j++) {
37  fr_t x1;
38  fr_cpy(x1, x2x4);
39 
40  fr_x2(x2x4);
41  fr_x4(x2x4);
42 
43  fr_x8(x8);
44 
45  if (fr_neq(x2x4, x8)) {
46  pass = false;
47 
48  printf("%d: FAILED: inconsistent result\n", j);
49  fr_print("1 : ", x1);
50  fr_print("2*4 : ", x2x4);
51  fr_print("8 : ", x8);
52  }
53  ++count;
54  }
55  }
56 
57  // 2*2*2*2*2*2 == 4*4*4 == 8*8
58 
59  for (int i=0; pass && i<TESTVALS; i++) {
60  fr_t x2, x4, x8;
61 
62  fr_cpy(x2, testval[i]);
63  fr_cpy(x4, testval[i]);
64  fr_cpy(x8, testval[i]);
65 
66  for (int j=0; pass && j<ITER; j++) {
67  fr_t x1;
68  fr_cpy(x1, x2);
69 
70  fr_x2(x2);
71  fr_x2(x2);
72  fr_x2(x2);
73  fr_x2(x2);
74  fr_x2(x2);
75  fr_x2(x2);
76 
77  fr_x4(x4);
78  fr_x4(x4);
79  fr_x4(x4);
80 
81  fr_x8(x8);
82  fr_x8(x8);
83 
84  if (fr_neq(x2, x4) || fr_neq(x2, x8)) {
85  pass = false;
86 
87  printf("%d: FAILED: inconsistent result\n", j);
88  fr_print("1 : ", x1);
89  fr_print("2^6 : ", x2);
90  fr_print("4^3 : ", x4);
91  fr_print("8^2 : ", x8);
92  }
93  ++count;
94  }
95  }
96 
97  // 3*4 == 12
98 
99  for (int i=0; pass && i<TESTVALS; i++) {
100  fr_t x3x4, x12;
101 
102  fr_cpy(x3x4, testval[i]);
103  fr_cpy(x12, testval[i]);
104 
105  for (int j=0; pass && j<ITER; j++) {
106  fr_t x1;
107  fr_cpy(x1, x3x4);
108 
109  fr_x3(x3x4);
110  fr_x4(x3x4);
111 
112  fr_x12(x12);
113 
114  if (fr_neq(x3x4, x12)) {
115  pass = false;
116 
117  printf("%d: FAILED: inconsistent result\n", j);
118  fr_print("1 : ", x1);
119  fr_print("3*4 : ", x3x4);
120  fr_print("12 : ", x12);
121  }
122  ++count;
123  }
124  }
125 
126  // 12+8 == 4(3+2)
127 
128  for (int i=0; pass && i<TESTVALS; i++) {
129  fr_t x1, x2, x3, x8, x12, l, r;
130 
131  fr_cpy(l, testval[i]);
132  fr_cpy(r, testval[i]);
133 
134  for (int j=0; pass && j<ITER; j++) {
135 
136  fr_cpy(x1, l);
137 
138  fr_cpy(x2, l);
139  fr_cpy(x3, l);
140  fr_cpy(x8, l);
141  fr_cpy(x12, l);
142 
143  fr_x2(x2);
144  fr_x3(x3);
145  fr_x8(x8);
146  fr_x12(x12);
147 
148  fr_cpy(l, x12);
149  fr_add(l, x8);
150 
151  fr_cpy(r, x3);
152  fr_add(r, x2);
153  fr_x4(r);
154 
155  if (fr_neq(l, r)) {
156  pass = false;
157 
158  printf("%d: FAILED: inconsistent result\n", i);
159  fr_print("1 : ", x1);
160  fr_print("12+8 : ", l);
161  fr_print("4(3+2) : ", r);
162  }
163  ++count;
164  }
165  }
166 
167  // 3*3*3*2*4*8 == 12*12*12
168 
169  for (int i=0; pass && i<TESTVALS; i++) {
170  fr_t x1, l, r;
171 
172  fr_cpy(l, testval[i]);
173  fr_cpy(r, testval[i]);
174 
175  for (int j=0; pass && j<ITER; j++) {
176 
177  fr_cpy(x1, l);
178 
179  fr_x3(l);
180  fr_x3(l);
181  fr_x3(l);
182  fr_x2(l);
183  fr_x4(l);
184  fr_x8(l);
185 
186  fr_x12(r);
187  fr_x12(r);
188  fr_x12(r);
189 
190  if (fr_neq(l, r)) {
191  pass = false;
192 
193  printf("%d: FAILED: inconsistent result\n", i);
194  fr_print("1 : ", x1);
195  fr_print("3*3*3*2*4*8 : ", l);
196  fr_print("12*12*12 : ", r);
197  }
198  ++count;
199  }
200  }
201 
202  printf("%ld tests\n", count);
203 
204  PRINTPASS(pass);
205 }
206 
207 // vim: ts=4 et sw=4 si
__managed__ testval_t testval[TESTVALS]
Definition: fptest.cu:8
#define TESTVALS
Definition: fptest.cuh:13
__device__ void fr_print(const char *s, const fr_t &x)
prints the canonical representation of x to STDOUT.
Definition: fr.cu:41
__device__ void fr_x12(fr_t &z)
Multiply the residue mod r z by 12 with weak reduction.
Definition: fr_x12.cu:13
__device__ void fr_x4(fr_t &z)
Multiply z by 4, and stores in z, with weak reduction.
Definition: fr_x4.cu:13
__device__ void fr_x2(fr_t &z)
Multiply z by 2, and stores in z, with weak reduction.
Definition: fr_x2.cu:13
__device__ void fr_x8(fr_t &z)
Multiply z by 8, and stores in z, with weak reduction.
Definition: fr_x8.cu:13
__device__ bool fr_neq(const fr_t &x, const fr_t &y)
Compares two fr_t residues.
Definition: fr_neq.cu:15
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
__device__ void fr_x3(fr_t &z)
Multiply z by 3, and stores in z, with weak reduction.
Definition: fr_x3.cu:12
__device__ __host__ void fr_cpy(fr_t &z, const fr_t &x)
Copy from x into z.
Definition: fr_cpy.cu:14
__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
__global__ void FrTestMulConst(testval_t *testval)
Test multiply by constant on Fr:
#define ITER
#define PRINTPASS(pass)
Definition: test.h:25