FK20 CUDA
frtest_addsub.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 5
9 
21 __global__ void FrTestAddSub(testval_t *testval) {
22 
23  printf("=== RUN %s\n", __func__);
24 
25  bool pass = true;
26  size_t count = 0;
27 
28  for (int i=0; i<TESTVALS; i++) {
29  for (int j=0; j<TESTVALS; j++) {
30  fr_t a, b, x, y;
31 
32  fr_cpy(a, testval[i]);
33  fr_cpy(b, testval[i]);
34 
35  fr_cpy(x, testval[i]);
36  fr_cpy(y, testval[j]);
37 
38  fr_addsub(x, y); // (x,y) = (x+y,x-y)
39 
40  fr_add(a, testval[j]); // x + y
41  fr_sub(b, testval[j]); // x - y
42 
43  if (fr_neq(a, x) || fr_neq(b, y)) {
44  pass = false;
45 
46  printf("%d,%d: FAILED: inconsistent result\n", i, j);
47  fr_print("x = ", testval[i]);
48  fr_print("y = ", testval[j]);
49  fr_print("x+y = ", a);
50  fr_print("x+y = ", x);
51  fr_print("x-y = ", b);
52  fr_print("x-y = ", y);
53  }
54  ++count;
55  }
56  }
57 
58  for (int i=0; i<TESTVALS; i++) {
59  for (int j=0; j<TESTVALS; j++) {
60  fr_t a, b, x, y;
61 
62  fr_cpy(a, testval[i]);
63  fr_cpy(b, testval[j]);
64 
65  fr_cpy(x, testval[i]);
66  fr_cpy(y, testval[j]);
67 
68  for (int k=0; k<ITER; k++) {
69 
70  // a,b -> 2a, 2b
71  fr_x2(a);
72  fr_x2(b);
73 
74  // x,y -> 2x, 2y
75  fr_addsub(x, y);
76  fr_addsub(x, y);
77 
78  if (fr_neq(a, x) || fr_neq(b, y)) {
79  pass = false;
80 
81  printf("%d,%d,%d: FAILED: inconsistent result\n", i, j, k);
82  printf("[%d]: ", i); fr_print("", testval[i]);
83  printf("[%d]: ", j); fr_print("", testval[j]);
84  fr_print("a = ", a);
85  fr_print("x = ", x);
86  fr_print("b = ", b);
87  fr_print("y = ", y);
88  }
89  ++count;
90  }
91  }
92  }
93  printf("%ld tests\n", count);
94 
95  PRINTPASS(pass);
96 }
97 
98 // 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_sub(fr_t &z, const fr_t &x)
Calculates the difference of two residues modulo p and stores it into z.
Definition: fr_sub.cu:17
__device__ void fr_x2(fr_t &z)
Multiply z by 2, and stores in z, with weak reduction.
Definition: fr_x2.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_addsub(fr_t &x, fr_t &y)
Computes the sum and the difference of the arguments, storing back into the arguments: (x,...
Definition: fr_addsub.cu:18
__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 FrTestAddSub(testval_t *testval)
Test for the fr_addsub kernel.
#define ITER
Definition: frtest_addsub.cu:8
#define PRINTPASS(pass)
Definition: test.h:25