Skip to content

Commit c3fdeae

Browse files
author
Luan Cardoso
committed
Fib Tests PTX
1 parent 0458790 commit c3fdeae

6 files changed

+334
-46
lines changed

‎Makefile

+1-1
Original file line numberDiff line numberDiff line change
@@ -10,7 +10,7 @@ FR=fr fr_cpy fr_reduce4 fr_eq fr_neq fr_neg fr_x2 fr_x3 fr_x4 fr_x8 fr_x12 fr_ad
1010
G1=g1a g1p g1p_compare g1p_add g1p_dbl g1p_mul g1p_neg g1p_scale g1p_ispoint g1p_sub g1p_addsub g1p_fft g1p_ptx
1111
FK20=fk20 fk20_poly2h_fft fk20_poly2toeplitz_coefficients fk20_poly2toeplitz_coefficients_fft fk20_poly2hext_fft fk20_msm fk20_hext_fft2h_fft
1212

13-
FPTEST=test fptest fptest_kat fptest_cmp fptest_mma fptest_inv fptest_add fptest_sub fptest_mul fptest_mulconst fptest_sqr fptest_distributive fptest_fibonacci fp_ptx
13+
FPTEST=test fptest fptest_kat fptest_cmp fptest_mma fptest_inv fptest_add fptest_sub fptest_mul fptest_mulconst fptest_sqr fptest_distributive fptest_fibonacci fp_ptx fpPTXtest
1414
FRTEST=test frtest frtest_kat frtest_cmp frtest_add frtest_mul frtest_inv frtest_sub frtest_addsub frtest_fibonacci frtest_mulconst frtest_sqr frtest_distributive frtest_fft
1515
G1TEST=test g1test g1test_kat g1test_fibonacci g1test_dbl g1test_fft
1616
FK20TEST=test fk20test fk20test_poly fk20_testvector fk20test_fft fk20test_fft_rand

‎fpPTXtest.cu

+205
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,205 @@
1+
// bls12_381: Arithmetic for BLS12-381
2+
// Copyright 2022-2023 Dag Arne Osvik
3+
// Copyright 2022-2023 Luan Cardoso dos Santos
4+
5+
6+
// Tests if the inline PTX and pure PTX functions generate matching results.
7+
8+
#include "fp.cuh"
9+
#include "fptest.cuh"
10+
#include "fp_ptx.cuh"
11+
12+
/**
13+
* @brief Test the equivalence between the inline PTX functions with the pure PTX
14+
* macros
15+
*
16+
* @return __global__
17+
*/
18+
__global__ void FpTestEqPTXInline(testval_t *testval) {
19+
printf("=== RUN %s\n", __func__);
20+
21+
bool pass = true;
22+
size_t count = 0;
23+
fp_t x, y, r, rp;
24+
//add
25+
for(int i=0; pass && i<TESTVALS; i++){
26+
for(int j=0; pass && j<TESTVALS; j++){
27+
fp_cpy(x, testval[i]);
28+
fp_cpy(y, testval[j]);
29+
30+
fp_add(r, x, y);
31+
fp_add_ptx(rp, x, y);
32+
33+
if(fp_neq(r, rp)){
34+
pass = false;
35+
36+
printf("%d: FAILED add\n", i);
37+
fp_print("x : ", x);
38+
fp_print("y : ", y);
39+
fp_print("r : ", r);
40+
fp_print("rp : ", rp);
41+
}
42+
++count;
43+
}
44+
}
45+
//sub
46+
for(int i=0; pass && i<TESTVALS; i++){
47+
for(int j=0; pass && j<TESTVALS; j++){
48+
fp_cpy(x, testval[i]);
49+
fp_cpy(y, testval[j]);
50+
51+
fp_sub(r, x, y);
52+
fp_sub_ptx(rp, x, y);
53+
54+
if(fp_neq(r, rp)){
55+
pass = false;
56+
57+
printf("%d: FAILED sub\n", i);
58+
fp_print("x : ", x);
59+
fp_print("y : ", y);
60+
fp_print("r : ", r);
61+
fp_print("rp : ", rp);
62+
}
63+
++count;
64+
}
65+
}
66+
//mul
67+
for(int i=0; pass && i<TESTVALS; i++){
68+
for(int j=0; pass && j<TESTVALS; j++){
69+
fp_cpy(x, testval[i]);
70+
fp_cpy(y, testval[j]);
71+
72+
fp_mul(r, x, y);
73+
fp_mul_ptx(rp, x, y);
74+
75+
if(fp_neq(r, rp)){
76+
pass = false;
77+
78+
printf("%d: FAILED mul\n", i);
79+
fp_print("x : ", x);
80+
fp_print("y : ", y);
81+
fp_print("r : ", r);
82+
fp_print("rp : ", rp);
83+
}
84+
++count;
85+
}
86+
}
87+
88+
//sqr
89+
for(int i=0; pass && i<TESTVALS; i++){
90+
fp_cpy(x, testval[i]);
91+
92+
fp_sqr(r, x);
93+
fp_sqr_ptx(rp, x);
94+
95+
if(fp_neq(r, rp)){
96+
pass = false;
97+
98+
printf("%d: FAILED sqr\n", i);
99+
fp_print("x : ", x);
100+
fp_print("y : ", y);
101+
fp_print("r : ", r);
102+
fp_print("rp : ", rp);
103+
}
104+
++count;
105+
}
106+
107+
//x2
108+
for(int i=0; pass && i<TESTVALS; i++){
109+
fp_cpy(x, testval[i]);
110+
111+
fp_x2(r, x);
112+
fp_x2_ptx(rp, x);
113+
114+
if(fp_neq(r, rp)){
115+
pass = false;
116+
117+
printf("%d: FAILED x2\n", i);
118+
fp_print("x : ", x);
119+
fp_print("y : ", y);
120+
fp_print("r : ", r);
121+
fp_print("rp : ", rp);
122+
}
123+
++count;
124+
}
125+
126+
//x3
127+
for(int i=0; pass && i<TESTVALS; i++){
128+
fp_cpy(x, testval[i]);
129+
130+
fp_x3(r, x);
131+
fp_x3_ptx(rp, x);
132+
133+
if(fp_neq(r, rp)){
134+
pass = false;
135+
136+
printf("%d: FAILED x2\n", i);
137+
fp_print("x : ", x);
138+
fp_print("y : ", y);
139+
fp_print("r : ", r);
140+
fp_print("rp : ", rp);
141+
}
142+
++count;
143+
}
144+
145+
//x4
146+
for(int i=0; pass && i<TESTVALS; i++){
147+
fp_cpy(x, testval[i]);
148+
149+
fp_x4(r, x);
150+
fp_x4_ptx(rp, x);
151+
152+
if(fp_neq(r, rp)){
153+
pass = false;
154+
155+
printf("%d: FAILED x2\n", i);
156+
fp_print("x : ", x);
157+
fp_print("y : ", y);
158+
fp_print("r : ", r);
159+
fp_print("rp : ", rp);
160+
}
161+
++count;
162+
}
163+
164+
//x8
165+
for(int i=0; pass && i<TESTVALS; i++){
166+
fp_cpy(x, testval[i]);
167+
168+
fp_x8(r, x);
169+
fp_x8_ptx(rp, x);
170+
171+
if(fp_neq(r, rp)){
172+
pass = false;
173+
174+
printf("%d: FAILED x2\n", i);
175+
fp_print("x : ", x);
176+
fp_print("y : ", y);
177+
fp_print("r : ", r);
178+
fp_print("rp : ", rp);
179+
}
180+
++count;
181+
}
182+
183+
//x12
184+
for(int i=0; pass && i<TESTVALS; i++){
185+
fp_cpy(x, testval[i]);
186+
187+
fp_x12(r, x);
188+
fp_x12_ptx(rp, x);
189+
190+
if(fp_neq(r, rp)){
191+
pass = false;
192+
193+
printf("%d: FAILED x2\n", i);
194+
fp_print("x : ", x);
195+
fp_print("y : ", y);
196+
fp_print("r : ", r);
197+
fp_print("rp : ", rp);
198+
}
199+
++count;
200+
}
201+
202+
printf("%ld tests\n", count);
203+
204+
PRINTPASS(pass);
205+
}

‎fp_ptx.ptxm

+2-7
Original file line numberDiff line numberDiff line change
@@ -17,11 +17,11 @@
1717
#include "fp_mul.ptxh"
1818
#include "fp_reduce12.ptxh"
1919
#include "fp_sqr.ptxh"
20-
#include "fp_x12.ptxh"
2120
#include "fp_x2.ptxh"
2221
#include "fp_x4.ptxh"
2322
#include "fp_x3.ptxh"
2423
#include "fp_x8.ptxh"
24+
#include "fp_x12.ptxh"
2525

2626

2727
// .globl _Z10fp_add_ptxRA6_mRA6_KmS3_
@@ -89,7 +89,6 @@
8989

9090
}
9191

92-
9392
// .globl _Z10fp_sub_ptxRA6_mRA6_KmS3_
9493
.visible .func _Z10fp_sub_ptxRA6_mRA6_KmS3_(
9594
.param .b64 _Z6fp_subRA6_mRA6_KmS3__param_0,
@@ -155,7 +154,6 @@
155154
}
156155

157156
// .globl _Z10fp_mul_ptxRA6_mRA6_KmS3_
158-
159157
.visible .func _Z10fp_mul_ptxRA6_mRA6_KmS3_(
160158
.param .b64 _Z10fp_mul_ptxRA6_mRA6_KmS3__param_0,
161159
.param .b64 _Z10fp_mul_ptxRA6_mRA6_KmS3__param_1,
@@ -221,9 +219,7 @@
221219

222220
}
223221

224-
225222
// .globl _Z10fp_sqr_ptxRA6_mRA6_Km
226-
227223
.visible .func _Z10fp_sqr_ptxRA6_mRA6_Km(
228224
.param .b64 _Z10fp_sqr_ptxRA6_mRA6_Km_param_0,
229225
.param .b64 _Z10fp_sqr_ptxRA6_mRA6_Km_param_1
@@ -275,7 +271,6 @@
275271

276272
}
277273

278-
279274
// .globl _Z9fp_x2_ptxRA6_mRA6_Km
280275
.visible .func _Z9fp_x2_ptxRA6_mRA6_Km(
281276
.param .b64 _Z9fp_x2_ptxRA6_mRA6_Km_param_0,
@@ -327,7 +322,6 @@
327322
}
328323

329324
// .globl _Z9fp_x3_ptxRA6_mRA6_Km
330-
331325
.visible .func _Z9fp_x3_ptxRA6_mRA6_Km(
332326
.param .b64 _Z9fp_x3_ptxRA6_mRA6_Km_param_0,
333327
.param .b64 _Z9fp_x3_ptxRA6_mRA6_Km_param_1
@@ -376,6 +370,7 @@
376370
ret;
377371

378372
}
373+
379374
// .globl _Z9fp_x4_ptxRA6_mRA6_Km
380375
.visible .func _Z9fp_x4_ptxRA6_mRA6_Km(
381376
.param .b64 _Z9fp_x4_ptxRA6_mRA6_Km_param_0,

‎fptest.cu

+2
Original file line numberDiff line numberDiff line change
@@ -124,6 +124,8 @@ int main(int argc, char **argv) {
124124
}
125125

126126
TEST(FpTestFibonacci);
127+
TEST(FpTestFibonacciPTX);
128+
TEST(FpTestEqPTXInline);
127129

128130
if (level >= 1) {
129131
printf("=== Tests level 1\n");

‎fptest.cuh

+2
Original file line numberDiff line numberDiff line change
@@ -25,6 +25,8 @@ typedef fp_t testval_t;
2525

2626
TESTFUN(FpTestKAT);
2727
TESTFUN(FpTestFibonacci);
28+
TESTFUN(FpTestFibonacciPTX);
29+
TESTFUN(FpTestEqPTXInline);
2830

2931
// Linear
3032

0 commit comments

Comments
 (0)