FK20 CUDA
fptest_mul.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 #include "fptest.cuh"
7 
17 
18  printf("=== RUN %s\n", __func__);
19 
20  bool pass = true;
21  size_t count = 0;
22 
23  for (int i=0; i<TESTVALS; i++) {
24  for (int j=0; j<TESTVALS; j++) {
25  fp_t x, y;
26 
27  fp_cpy(x, testval[i]);
28  fp_cpy(y, testval[j]);
29 
30  fp_mul(x, x, testval[j]); // x * y
31  fp_mul(y, y, testval[i]); // y * x
32 
33  if (fp_neq(x, y)) {
34  pass = false;
35 
36  printf("%d,%d: FAILED: inconsistent result\n", i, j);
37  fp_print("x = ", testval[i]);
38  fp_print("y = ", testval[j]);
39  fp_print("x*y = ", x);
40  fp_print("y*x = ", y);
41  }
42  ++count;
43  }
44  }
45  printf("%ld tests\n", count);
46 
47  PRINTPASS(pass);
48 }
49 
50 //
51 
61 
62  printf("=== RUN %s\n", __func__);
63 
64  bool pass = true;
65  size_t count = 0;
66 
67  for (int i=0; i<TESTVALS; i++) {
68  for (int j=0; j<TESTVALS; j++) {
69  for (int k=0; k<TESTVALS; k++) {
70  fp_t a, b, c;
71 
72  fp_cpy(a, testval[i]); // x
73  fp_cpy(b, testval[j]); // y
74  fp_cpy(c, testval[i]); // x
75 
76  fp_mul(a, a, testval[j]); // x * y
77  fp_mul(a, a, testval[k]); // (x * y) * z
78 
79  fp_mul(b, b, testval[k]); // y * z
80  fp_mul(c, c, b); // x * (y * z)
81 
82  if (fp_neq(a, c)) {
83  pass = false;
84 
85  printf("%d,%d,%d: FAILED: inconsistent result\n", i, j, k);
86  fp_print("x = ", testval[i]);
87  fp_print("y = ", testval[j]);
88  fp_print("z = ", testval[k]);
89  fp_print("(x*y)*z = ", a);
90  fp_print("x*(y*z) = ", c);
91  }
92  ++count;
93  }
94  }
95  }
96  printf("%ld tests\n", count);
97 
98  PRINTPASS(pass);
99 }
100 
109 __global__ void FpTestMul(testval_t *testval) {
110 
111  printf("=== RUN %s\n", __func__);
112 
113  bool pass = true;
114 
115  unsigned tid = 0; tid += blockIdx.z;
116  tid *= gridDim.y; tid += blockIdx.y;
117  tid *= gridDim.x; tid += blockIdx.x;
118  tid *= blockDim.z; tid += threadIdx.z;
119  tid *= blockDim.y; tid += threadIdx.y;
120  tid *= blockDim.x; tid += threadIdx.x;
121 
122  if (tid > 0xffffffff)
123  return;
124 
125  uint64_t t[6] = { tid, 0, 0, 0, 0, 0 };
126 
127  fp_t x, z;
128 
129  fp_fromUint64(x, t); // x = tid
130 
131  if (x[0] != tid) pass = false;
132  if (x[1] != 0) pass = false;
133  if (x[2] != 0) pass = false;
134  if (x[3] != 0) pass = false;
135  if (x[4] != 0) pass = false;
136  if (x[5] != 0) pass = false;
137 
138  if (!pass) {
139  printf("%d: FAILED after fp_fromUint64\n", tid);
140  goto done;
141  }
142 
143  fp_mul(z, x, x); // z = tid * tid
144 
145  if (z[0] != tid*tid) pass = false;
146  if (z[1] != 0) pass = false;
147  if (z[2] != 0) pass = false;
148  if (z[3] != 0) pass = false;
149  if (z[4] != 0) pass = false;
150  if (z[5] != 0) pass = false;
151 
152  if (!pass) {
153  printf("%d: FAILED after fp_mul\n", tid);
154  goto done;
155  }
156 
157 done:
158  PRINTPASS(pass);
159 }
160 
161 // vim: ts=4 et sw=4 si
__device__ void fp_print(const char *s, const fp_t &x)
Prints the canonical representation of x to STDOUT.
Definition: fp.cu:39
__device__ __host__ void fp_fromUint64(fp_t &z, const uint64_t *x)
Converts uint64_t[6] to fp_t. After this operation, z represents x mod p.
Definition: fp.cu:58
__device__ bool fp_neq(const fp_t &x, const fp_t &y)
Compares two fp_t residues.
Definition: fp_neq.cu:14
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
__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
__device__ __host__ void fp_cpy(fp_t &z, const fp_t &x)
Copy from x into z.
Definition: fp_cpy.cu:14
__managed__ testval_t testval[TESTVALS]
Definition: fptest.cu:8
#define TESTVALS
Definition: fptest.cuh:13
__global__ void FpTestMul(testval_t *testval)
Multiplication test, comparing with the native uint64_t multiplication.
Definition: fptest_mul.cu:109
__global__ void FpTestAssociativeMul(testval_t *testval)
Test for the associative property of multiplication.
Definition: fptest_mul.cu:60
__global__ void FpTestCommutativeMul(testval_t *testval)
Test for the commutative property of addition.
Definition: fptest_mul.cu:16
#define PRINTPASS(pass)
Definition: test.h:25