FK20 CUDA
fptest_inv.cu
Go to the documentation of this file.
1 // bls12_381: Arithmetic for BLS12-381
2 // Copyright 2022 Dag Arne Osvik
3 
4 #include "fp.cuh"
5 #include "fptest.cuh"
6 
15 __global__ void FpTestInv(testval_t *testval) {
16 
17  printf("=== RUN %s\n", __func__);
18 
19  bool pass = true;
20 
21  unsigned tid = 0; tid += blockIdx.z;
22  tid *= gridDim.y; tid += blockIdx.y;
23  tid *= gridDim.x; tid += blockIdx.x;
24  tid *= blockDim.z; tid += threadIdx.z;
25  tid *= blockDim.y; tid += threadIdx.y;
26  tid *= blockDim.x; tid += threadIdx.x;
27 
28  uint64_t v[6] = { tid + 2, 0, 0, 0, 0, 0 };
29 
30  fp_t x, y, z;
31 
32  fp_fromUint64(x, v); // x = tid + 2
33 
34  if (x[0] != tid+2) pass = false;
35  if (x[1] != 0) pass = false;
36  if (x[2] != 0) pass = false;
37  if (x[3] != 0) pass = false;
38  if (x[4] != 0) pass = false;
39  if (x[5] != 0) pass = false;
40 
41  if (!pass) {
42  printf("%d: FAILED after fp_fromUint64\n", tid);
43  goto done;
44  }
45 
46  fp_inv(y, x); // y = x^-1
47 
48  if (y[0] == tid+2) pass = false;
49 
50  if (!pass) {
51  printf("%d: FAILED after fp_inv\n", tid);
52  goto done;
53  }
54 
55  fp_mul(z, y, x); // z = y * x
56 
57  fp_reduce6(z);
58 
59  if (z[0] != 1) pass = false;
60  if (z[1] != 0) pass = false;
61  if (z[2] != 0) pass = false;
62  if (z[3] != 0) pass = false;
63  if (z[4] != 0) pass = false;
64  if (z[5] != 0) pass = false;
65 
66  if (!pass) {
67  printf("%d: FAILED test of inverse\n", tid);
68  goto done;
69  }
70 
71  fp_mul(z, z, x); // z *= x
72 
73  fp_reduce6(z);
74 
75  if (z[0] != v[0]) pass = false;
76  if (z[1] != v[1]) pass = false;
77  if (z[2] != v[2]) pass = false;
78  if (z[3] != v[3]) pass = false;
79  if (z[4] != v[4]) pass = false;
80  if (z[5] != v[5]) pass = false;
81 
82  if (!pass) {
83  printf("%d: FAILED to match input value\n", tid);
84  goto done;
85  }
86 
87 done:
88  PRINTPASS(pass);
89 }
90 
91 // vim: ts=4 et sw=4 si
__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__ void fp_reduce6(fp_t &z)
Narrow reduction of a residue modulo p, reducing to the canonical representation.
Definition: fp_reduce6.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__ 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
__managed__ testval_t testval[TESTVALS]
Definition: fptest.cu:8
__global__ void FpTestInv(testval_t *testval)
Test for multiplicative inverse mod p in Fp.
Definition: fptest_inv.cu:15
#define PRINTPASS(pass)
Definition: test.h:25