From 200642b99fc3c78818938109507da0786aaefd9b Mon Sep 17 00:00:00 2001 From: Craig Raw Date: Wed, 15 Oct 2025 13:44:34 +0200 Subject: [PATCH] adapt curve from sm2 to secp256k1 --- .DS_Store | Bin 6148 -> 0 bytes .gitignore | 3 ++- include/.DS_Store | Bin 6148 -> 0 bytes include/gecc/.DS_Store | Bin 6148 -> 0 bytes include/gecc/arith/details/fp_mont_multiply.h | 4 ++-- include/gecc/arith/fp.h | 4 ++-- scripts/constants.py | 19 +++++++++++------- scripts/constants_generator.py | 16 +++++++-------- scripts/ec.py | 8 ++++---- scripts/field.py | 6 +++--- test/ecdsa_ec_fixed_pmul.cu | 8 ++++---- test/ecdsa_ec_unknown_pmul.cu | 8 ++++---- test/ecdsa_sign.cu | 8 ++++---- test/ecdsa_sign_baseline.cu | 10 ++++----- test/ecdsa_verify.cu | 8 ++++---- test/ecdsa_verify_baseline.cu | 8 ++++---- test/fp.cu | 16 +++++++-------- test/modinv_data_parallel_profiling.cu | 8 ++++---- 18 files changed, 70 insertions(+), 64 deletions(-) delete mode 100644 .DS_Store delete mode 100644 include/.DS_Store delete mode 100644 include/gecc/.DS_Store diff --git a/.DS_Store b/.DS_Store deleted file mode 100644 index d11bd01539510182b0e1f2ebb639b08064e6c0af..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 6148 zcmeHKu}<7T5Pb_Lf}nr`MG=)NI&{#{AQ|K+DI%IvsKMNEAkJqC`ywTkl`=n&7DdW0 zq)QQj@&o(?pTV2m1-v$nl;#9z2knkM`)23)*|*CAP`i`!sN|Ve`Ui#UNOwuTd zM`=uiEELDrtCu_fLihgzot;!!gq#+yaEdXGFyd-M<_=HQ``biR`Gw!MzF)6h3BTGr zN&3L69p&(uIp8#$CnuY8vdT~(clhAfO%qY&!_hmvYV%b2w4)q8T!m>JvKWw4?G66^ zQ5|dHaG%>!{eIl7yCHSw_zeyCX8N-JS5;8)*OmHsOMPdK77h-k%0vCRXz<92_`@VG zWIx(Q51&OBvP{s$K6%PFlxVfOh;YcULymkTYQQ&|^I9Tf3~@r_Dc`eri3mRBy~!UF zC+K&?6nM{Q%)24J;EZ2IxBm2GwP#s5Pb2iFX69?b{NmIq`a)g-ufTIrK-PzVO?+s5 zX}n%K*y$qxeTQLfyq2?sCgTm@#$?Xv9uZ`;rRZj!(A3U~!xlme;~^`c!qlHFSy4<~zV uz;?hUCb-ggox;v;$NECH;x}yCI2OnSeQ14YJR>y!kARlJcV2<1DsT&vD$V8q diff --git a/.gitignore b/.gitignore index 863fbc0..1fcc698 100644 --- a/.gitignore +++ b/.gitignore @@ -17,4 +17,5 @@ __pycache__/ *$py.class *.sage.py -.ipynb_checkpoints/ \ No newline at end of file +.ipynb_checkpoints/ +.DS_Store diff --git a/include/.DS_Store b/include/.DS_Store deleted file mode 100644 index ed7ee9aca284dcd48cdaeaed92c44f97f79e870e..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 6148 zcmeHKJ5Iw;5ZnzeBGIHoRPP9+rg0*3f?NP4Q9vS*)1;wt$H8)hT!`5Z$d(mK3YeAl zjoAkPlZ;eAwR`Me=BrJ6IF2f>2aH?EMo#iB zZv(zQ?)tey+cfKayGG0&W_P(iFNXQ;kNM+k@7X2~pwg%S6`%rCfC~Ie1<%=b5KOIZG%D~L3VZ^p CHzue6 diff --git a/include/gecc/.DS_Store b/include/gecc/.DS_Store deleted file mode 100644 index ca4c0ba3c92cb3421561f8d97e9e04bde8ce4f3e..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 6148 zcmeHK%}&EG40cA5hS*_8#3^@9dxJ2QC)fi3+kj1K)$%j#THc69K;pm|@j`q~tZG@8 zaX>-{*^=|6j@`Ilqq-&{GhWR{L_;F#Lm7KV7mI2WCdn}5((PUPYb;nl? zf0F@zcZYOM545Bw+IoNM_jt9e@_brUQ&{;x+0zwxLSn98gDPKp9X5S_b(0V4;jounFit9TEm0_vy?C<8GA2jOza z_5UK>|Hnakr3@$o|B3;V<&%7jN0Qpwc{r}M7J3C`VZSC|8-j^1#qi})d;;|XyX6Vc U2{r*?f#{Ebr$Gm0;8z*=0{h@;AOHXW diff --git a/include/gecc/arith/details/fp_mont_multiply.h b/include/gecc/arith/details/fp_mont_multiply.h index 3d7dc65..abdb709 100644 --- a/include/gecc/arith/details/fp_mont_multiply.h +++ b/include/gecc/arith/details/fp_mont_multiply.h @@ -306,7 +306,7 @@ __device__ __forceinline__ FpT operator*(const FpT &b) const { return mont_multiply_cios(b.digits); } else if (mont_flag == MONTFLAG::SOS) { - if (curve_flag == CURVEFLAG::SM2) + if (curve_flag == CURVEFLAG::SECP256K1) return mont_multiply_sos_sm2(b.digits); else return mont_multiply_sos(b.digits); @@ -319,7 +319,7 @@ __device__ __forceinline__ FpT mont_multiply(const Base *o, return mont_multiply_cios(o, stride); } else if (mont_flag == MONTFLAG::SOS) { - if (curve_flag == CURVEFLAG::SM2) + if (curve_flag == CURVEFLAG::SECP256K1) return mont_multiply_sos_sm2(o, stride); else return mont_multiply_sos(o, stride); diff --git a/include/gecc/arith/fp.h b/include/gecc/arith/fp.h index 8255bc6..dcf6b3e 100644 --- a/include/gecc/arith/fp.h +++ b/include/gecc/arith/fp.h @@ -13,7 +13,7 @@ using FP_NAME = \ FpT -#define DEFINE_SM2_FP(FP_NAME, FP_TYPE, DIGIT_TYPE, DIGIT_WIDTH, LAYOUT, LIMBS, MONT_FLAG, CURVE_FLAG) \ +#define DEFINE_SECP256K1_FP(FP_NAME, FP_TYPE, DIGIT_TYPE, DIGIT_WIDTH, LAYOUT, LIMBS, MONT_FLAG, CURVE_FLAG) \ using FP_NAME##Factory = \ gecc::arith::FpFactory, LAYOUT, LIMBS>; \ __device__ __constant__ FP_NAME##Factory::Constant FP_NAME##DCONST; \ @@ -48,7 +48,7 @@ enum MONTFLAG { enum CURVEFLAG { DEFAULT, - SM2, + SECP256K1, // Optimized for secp256k1 curve (formerly SM2) }; template struct FpFactory { diff --git a/scripts/constants.py b/scripts/constants.py index 9995e71..e2c59e6 100644 --- a/scripts/constants.py +++ b/scripts/constants.py @@ -1,11 +1,16 @@ -SM2_q = 0xFFFFFFFEFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF00000000FFFFFFFFFFFFFFFF -SM2_g1_a = -3 -SM2_g1_b = 0x28E9FA9E9D9F5E344D5A9E4BCF6509A7F39789F515AB8F92DDBCBD414D940E93 -SM2_g1_generator = ( - 0x32c4ae2c1f1981195f9904466a39c9948fe30bbff2660be1715a4589334c74c7, - 0xbc3736a2f4f6779c59bdcee36b692153d0a9877cc62a474002df32e52139f0a0 +# secp256k1 curve parameters +# Field prime (same as curve order field) +SECP256K1_q = 0xFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFEFFFFFC2F +# Curve equation: y^2 = x^3 + ax + b +SECP256K1_g1_a = 0 +SECP256K1_g1_b = 7 +# Generator point +SECP256K1_g1_generator = ( + 0x79BE667EF9DCBBAC55A06295CE870B07029BFCDB2DCE28D959F2815B16F81798, + 0x483ADA7726A3C4655DA4FBFC0E1108A8FD17B448A68554199C47D08FFB10D4B8 ) -SM2_n = 0xfffffffeffffffffffffffffffffffff7203df6b21c6052b53bbf40939d54123 +# Curve order (number of points) +SECP256K1_n = 0xFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFEBAAEDCE6AF48A03BBFD25E8CD0364141 ECDSA_Verify_g1_generator = ( # 0x74c1082d5c8bc69cdc64d13d27ea1cfcc7b9d732d13f7f31f9fab63fdcf02b78, diff --git a/scripts/constants_generator.py b/scripts/constants_generator.py index 75912a0..2d650bd 100644 --- a/scripts/constants_generator.py +++ b/scripts/constants_generator.py @@ -277,19 +277,19 @@ if __name__ == '__main__': with open(root / 'fp_constants.h', 'w') as f: crepr_64 = CRepr() crepr_64.width = 64 - f.write(crepr_64.fp_constant(field.Fq_SM2) + '\n') - f.write(crepr_64.fp_constant(field.Fq_SM2_n) + '\n') + f.write(crepr_64.fp_constant(field.Fq_SECP256K1) + '\n') + f.write(crepr_64.fp_constant(field.Fq_SECP256K1_n) + '\n') with open(root / 'ec_constants.h', 'w') as f: crepr_64 = CRepr() crepr_64.width = 64 - f.write(crepr_64.ec_constant(ec.G1_SM2) + '\n') + f.write(crepr_64.ec_constant(ec.G1_SECP256K1) + '\n') # f.write(crepr_64.ec_constant(ec.G1_ECDSA_VERIFY) + '\n') with open(root / 'ecdsa_constants.h', 'w') as f: crepr_64 = CRepr() crepr_64.width = 64 - f.write(crepr_64.ecdsa_constant(ec.G1_SM2) + '\n') + f.write(crepr_64.ecdsa_constant(ec.G1_SECP256K1) + '\n') # f.write(crepr_64.ecdsa_constant(ec.G1_ECDSA_VERIFY) + '\n') with open(root / 'fp_ops_cc_details.h', 'w') as f: @@ -348,12 +348,12 @@ if __name__ == '__main__': # tests with open(root / 'fp_test_constants.h', 'w') as f: - generate_fp_test(f, field.Fq_SM2.name, field.Fq_SM2, 6, field.Fq_SM2.width) - generate_fp_test(f, field.Fq_SM2_n.name, field.Fq_SM2_n, 6, field.Fq_SM2_n.width) + generate_fp_test(f, field.Fq_SECP256K1.name, field.Fq_SECP256K1, 6, field.Fq_SECP256K1.width) + generate_fp_test(f, field.Fq_SECP256K1_n.name, field.Fq_SECP256K1_n, 6, field.Fq_SECP256K1_n.width) with open(root / 'ecdsa_test_constants.h', 'w') as f: generate_ecdsa_test( - f, field.Fq_SM2_n, field.Fq_SM2_n.width) + f, field.Fq_SECP256K1_n, field.Fq_SECP256K1_n.width) # generate_ecdsa_test( - # f, field.Fq_SM2, ec.G1_SM2, field.Fq_SM2.width) + # f, field.Fq_SECP256K1, ec.G1_SECP256K1, field.Fq_SECP256K1.width) diff --git a/scripts/ec.py b/scripts/ec.py index 0c2a29f..1fca452 100644 --- a/scripts/ec.py +++ b/scripts/ec.py @@ -171,10 +171,10 @@ class EC: def to_mont(self, p): return (type(p))(map(self.field.to_mont, p)) -G1_SM2 = EC('G1SM2', field.Fq_SM2, - constants.SM2_g1_a, constants.SM2_g1_b, generator=constants.SM2_g1_generator) -G1_ECDSA_VERIFY = EC('G1ECDSA_VERIFY', field.Fq_SM2, - constants.SM2_g1_a, constants.SM2_g1_b, generator=constants.ECDSA_Verify_g1_generator) +G1_SECP256K1 = EC('G1SECP256K1', field.Fq_SECP256K1, + constants.SECP256K1_g1_a, constants.SECP256K1_g1_b, generator=constants.SECP256K1_g1_generator) +G1_ECDSA_VERIFY = EC('G1ECDSA_VERIFY', field.Fq_SECP256K1, + constants.SECP256K1_g1_a, constants.SECP256K1_g1_b, generator=constants.ECDSA_Verify_g1_generator) def test_ec(self, ec): affine_p = ec.random_element() diff --git a/scripts/field.py b/scripts/field.py index 7614030..8ae91cf 100644 --- a/scripts/field.py +++ b/scripts/field.py @@ -1,4 +1,4 @@ -from constants import SM2_q, SM2_n +from constants import SECP256K1_q, SECP256K1_n import random @@ -187,5 +187,5 @@ class Fp2: a, b = x return (self.fp.to_mont(a), self.fp.to_mont(b)) -Fq_SM2 = Fp('FqSM2', SM2_q, 64, rexp=4) -Fq_SM2_n = Fp('FqSM2_n', SM2_n, 64, rexp=4) \ No newline at end of file +Fq_SECP256K1 = Fp('FqSECP256K1', SECP256K1_q, 64, rexp=4) +Fq_SECP256K1_n = Fp('FqSECP256K1_n', SECP256K1_n, 64, rexp=4) \ No newline at end of file diff --git a/test/ecdsa_ec_fixed_pmul.cu b/test/ecdsa_ec_fixed_pmul.cu index 6f66251..78acac9 100644 --- a/test/ecdsa_ec_fixed_pmul.cu +++ b/test/ecdsa_ec_fixed_pmul.cu @@ -56,8 +56,8 @@ void test_ecdsa_ec_fixed_pmul() { } -DEFINE_SM2_FP(Fq_SM2_1, FqSM2, u32, 32, LayoutT<1>, 8, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SM2); -DEFINE_FP(Fq_SM2_n, FqSM2_n, u32, 32, LayoutT<1>, 8); -DEFINE_EC(G1_1, G1SM2, Fq_SM2_1, SM2_CURVE, 2); -DEFINE_ECDSA(ECDSA_EC_PMUL_Solver, G1_1_G1SM2, Fq_SM2_1, Fq_SM2_n); +DEFINE_SECP256K1_FP(Fq_SECP256K1_1, FqSECP256K1, u32, 32, LayoutT<1>, 8, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SECP256K1); +DEFINE_FP(Fq_SECP256K1_n, FqSECP256K1_n, u32, 32, LayoutT<1>, 8); +DEFINE_EC(G1_1, G1SECP256K1, Fq_SECP256K1_1, SECP256K1_CURVE, 2); +DEFINE_ECDSA(ECDSA_EC_PMUL_Solver, G1_1_G1SECP256K1, Fq_SECP256K1_1, Fq_SECP256K1_n); TEST(ECDSA_EC_PMUL, Performance) { test_ecdsa_ec_fixed_pmul(); } \ No newline at end of file diff --git a/test/ecdsa_ec_unknown_pmul.cu b/test/ecdsa_ec_unknown_pmul.cu index 3f19eee..6707dca 100644 --- a/test/ecdsa_ec_unknown_pmul.cu +++ b/test/ecdsa_ec_unknown_pmul.cu @@ -56,8 +56,8 @@ void test_ecdsa_ec_unknown_pmul() { } -DEFINE_SM2_FP(Fq_SM2_1, FqSM2, u32, 32, LayoutT<1>, 8, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SM2); -DEFINE_FP(Fq_SM2_n, FqSM2_n, u32, 32, LayoutT<1>, 8); -DEFINE_EC(G1_1, G1SM2, Fq_SM2_1, SM2_CURVE, 2); -DEFINE_ECDSA(ECDSA_EC_PMUL_Solver, G1_1_G1SM2, Fq_SM2_1, Fq_SM2_n); +DEFINE_SECP256K1_FP(Fq_SECP256K1_1, FqSECP256K1, u32, 32, LayoutT<1>, 8, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SECP256K1); +DEFINE_FP(Fq_SECP256K1_n, FqSECP256K1_n, u32, 32, LayoutT<1>, 8); +DEFINE_EC(G1_1, G1SECP256K1, Fq_SECP256K1_1, SECP256K1_CURVE, 2); +DEFINE_ECDSA(ECDSA_EC_PMUL_Solver, G1_1_G1SECP256K1, Fq_SECP256K1_1, Fq_SECP256K1_n); TEST(ECDSA_EC_PMUL, Performance) { test_ecdsa_ec_unknown_pmul(); } \ No newline at end of file diff --git a/test/ecdsa_sign.cu b/test/ecdsa_sign.cu index ea071f6..573c498 100644 --- a/test/ecdsa_sign.cu +++ b/test/ecdsa_sign.cu @@ -94,9 +94,9 @@ void test_ecdsa_sign() { } -DEFINE_SM2_FP(Fq_SM2_1, FqSM2, u32, 32, LayoutT<1>, 8, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SM2); -DEFINE_FP(Fq_SM2_n, FqSM2_n, u32, 32, LayoutT<1>, 8); -DEFINE_EC(G1_1, G1SM2, Fq_SM2_1, SM2_CURVE, 2); -DEFINE_ECDSA(ECDSA_solver, G1_1_G1SM2, Fq_SM2_1, Fq_SM2_n); +DEFINE_SECP256K1_FP(Fq_SECP256K1_1, FqSECP256K1, u32, 32, LayoutT<1>, 8, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SECP256K1); +DEFINE_FP(Fq_SECP256K1_n, FqSECP256K1_n, u32, 32, LayoutT<1>, 8); +DEFINE_EC(G1_1, G1SECP256K1, Fq_SECP256K1_1, SECP256K1_CURVE, 2); +DEFINE_ECDSA(ECDSA_solver, G1_1_G1SECP256K1, Fq_SECP256K1_1, Fq_SECP256K1_n); TEST(ECDSA, Correctness) { test_ecdsa_sign_correctness(); } TEST(ECDSA, Performance) { test_ecdsa_sign(); } \ No newline at end of file diff --git a/test/ecdsa_sign_baseline.cu b/test/ecdsa_sign_baseline.cu index 462f2bf..5967e6c 100644 --- a/test/ecdsa_sign_baseline.cu +++ b/test/ecdsa_sign_baseline.cu @@ -94,10 +94,10 @@ void test_ecdsa_sign() { // TODO OPT -DEFINE_SM2_FP(Fq_SM2_1, FqSM2, u32, 32, LayoutT<1>, 8, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SM2); -// DEFINE_FP(Fq_SM2_1, FqSM2, u32, 32, LayoutT<1>, 8); -DEFINE_FP(Fq_SM2_n, FqSM2_n, u32, 32, LayoutT<1>, 8); -DEFINE_EC(G1_1, G1SM2, Fq_SM2_1, SM2_CURVE, 2); -DEFINE_ECDSA(ECDSA_solver, G1_1_G1SM2, Fq_SM2_1, Fq_SM2_n); +DEFINE_SECP256K1_FP(Fq_SECP256K1_1, FqSECP256K1, u32, 32, LayoutT<1>, 8, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SECP256K1); +// DEFINE_FP(Fq_SECP256K1_1, FqSECP256K1, u32, 32, LayoutT<1>, 8); +DEFINE_FP(Fq_SECP256K1_n, FqSECP256K1_n, u32, 32, LayoutT<1>, 8); +DEFINE_EC(G1_1, G1SECP256K1, Fq_SECP256K1_1, SECP256K1_CURVE, 2); +DEFINE_ECDSA(ECDSA_solver, G1_1_G1SECP256K1, Fq_SECP256K1_1, Fq_SECP256K1_n); TEST(ECDSA, Correctness) { test_ecdsa_sign_correctness(); } TEST(ECDSA, Performance) { test_ecdsa_sign(); } \ No newline at end of file diff --git a/test/ecdsa_verify.cu b/test/ecdsa_verify.cu index 73a5993..4b383ae 100644 --- a/test/ecdsa_verify.cu +++ b/test/ecdsa_verify.cu @@ -105,9 +105,9 @@ void test_ecdsa_verify() { } -DEFINE_SM2_FP(Fq_SM2_1, FqSM2, u32, 32, LayoutT<1>, 8, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SM2); -DEFINE_FP(Fq_SM2_n, FqSM2_n, u32, 32, LayoutT<1>, 8); -DEFINE_EC(G1_1, G1SM2, Fq_SM2_1, SM2_CURVE, 2); -DEFINE_ECDSA(ECDSA_Verify_Solver, G1_1_G1SM2, Fq_SM2_1, Fq_SM2_n); +DEFINE_SECP256K1_FP(Fq_SECP256K1_1, FqSECP256K1, u32, 32, LayoutT<1>, 8, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SECP256K1); +DEFINE_FP(Fq_SECP256K1_n, FqSECP256K1_n, u32, 32, LayoutT<1>, 8); +DEFINE_EC(G1_1, G1SECP256K1, Fq_SECP256K1_1, SECP256K1_CURVE, 2); +DEFINE_ECDSA(ECDSA_Verify_Solver, G1_1_G1SECP256K1, Fq_SECP256K1_1, Fq_SECP256K1_n); TEST(ECDSA, Correctness) { test_ecdsa_verify_correctness(); } TEST(ECDSA, Performance) { test_ecdsa_verify(); } \ No newline at end of file diff --git a/test/ecdsa_verify_baseline.cu b/test/ecdsa_verify_baseline.cu index d53b976..ad98bce 100644 --- a/test/ecdsa_verify_baseline.cu +++ b/test/ecdsa_verify_baseline.cu @@ -99,9 +99,9 @@ void test_ecdsa_verify() { } // Baseline: OPT ModMUL + rapid_EC -DEFINE_SM2_FP(Fq_SM2_1, FqSM2, u32, 32, LayoutT<1>, 8, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SM2); -DEFINE_FP(Fq_SM2_n, FqSM2_n, u32, 32, LayoutT<1>, 8); -DEFINE_EC(G1_1, G1SM2, Fq_SM2_1, SM2_CURVE, 2); -DEFINE_ECDSA(ECDSA_Verify_Solver, G1_1_G1SM2, Fq_SM2_1, Fq_SM2_n); +DEFINE_SECP256K1_FP(Fq_SECP256K1_1, FqSECP256K1, u32, 32, LayoutT<1>, 8, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SECP256K1); +DEFINE_FP(Fq_SECP256K1_n, FqSECP256K1_n, u32, 32, LayoutT<1>, 8); +DEFINE_EC(G1_1, G1SECP256K1, Fq_SECP256K1_1, SECP256K1_CURVE, 2); +DEFINE_ECDSA(ECDSA_Verify_Solver, G1_1_G1SECP256K1, Fq_SECP256K1_1, Fq_SECP256K1_n); TEST(ECDSA, Correctness) { test_ecdsa_verify_correctness(); } TEST(ECDSA, Performance) { test_ecdsa_verify(); } diff --git a/test/fp.cu b/test/fp.cu index b69bdf3..77eed4a 100644 --- a/test/fp.cu +++ b/test/fp.cu @@ -155,14 +155,14 @@ template void test_fp(size_t N, cudaFree(out); } -#define ADD_FqSM2_FP_TEST(FIELD, DIGIT_TYPE, DIGIT_WIDTH, LAYOUT_WIDTH, LIMBS) \ - DEFINE_SM2_FP(FIELD, FqSM2, DIGIT_TYPE, DIGIT_WIDTH, LayoutT, LIMBS, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SM2); \ - TEST(FqSM2256K1_FP, FIELD##Correctness) { using namespace FqSM2_fp_test; test_fp(N, A, B, SUM, PROD); } +#define ADD_FqSECP256K1_FP_TEST(FIELD, DIGIT_TYPE, DIGIT_WIDTH, LAYOUT_WIDTH, LIMBS) \ + DEFINE_SECP256K1_FP(FIELD, FqSECP256K1, DIGIT_TYPE, DIGIT_WIDTH, LayoutT, LIMBS, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SECP256K1); \ + TEST(FqSECP256K1_FP, FIELD##Correctness) { using namespace FqSECP256K1_fp_test; test_fp(N, A, B, SUM, PROD); } -#define ADD_FqSM2_n_FP_TEST(FIELD, DIGIT_TYPE, DIGIT_WIDTH, LAYOUT_WIDTH, LIMBS) \ - DEFINE_FP(FIELD, FqSM2_n, DIGIT_TYPE, DIGIT_WIDTH, LayoutT, LIMBS); \ - TEST(FqSM2_FP_n, FIELD##Correctness) { using namespace FqSM2_n_fp_test; test_fp(N, A, B, SUM, PROD); } +#define ADD_FqSECP256K1_n_FP_TEST(FIELD, DIGIT_TYPE, DIGIT_WIDTH, LAYOUT_WIDTH, LIMBS) \ + DEFINE_FP(FIELD, FqSECP256K1_n, DIGIT_TYPE, DIGIT_WIDTH, LayoutT, LIMBS); \ + TEST(FqSECP256K1_FP_n, FIELD##Correctness) { using namespace FqSECP256K1_n_fp_test; test_fp(N, A, B, SUM, PROD); } -ADD_FqSM2_FP_TEST(Field_SM2, u32, 32, 1, 8) -ADD_FqSM2_n_FP_TEST(Field_SM2_n, u32, 32, 1, 8) \ No newline at end of file +ADD_FqSECP256K1_FP_TEST(Field_SECP256K1, u32, 32, 1, 8) +ADD_FqSECP256K1_n_FP_TEST(Field_SECP256K1_n, u32, 32, 1, 8) \ No newline at end of file diff --git a/test/modinv_data_parallel_profiling.cu b/test/modinv_data_parallel_profiling.cu index 6117329..93e101d 100644 --- a/test/modinv_data_parallel_profiling.cu +++ b/test/modinv_data_parallel_profiling.cu @@ -56,8 +56,8 @@ void test_modinv_in_data_parallel() { } -DEFINE_SM2_FP(Fq_SM2_1, FqSM2, u32, 32, LayoutT<1>, 8, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SM2); -DEFINE_FP(Fq_SM2_n, FqSM2_n, u32, 32, LayoutT<1>, 8); -DEFINE_EC(G1_1, G1SM2, Fq_SM2_1, SM2_CURVE, 2); -DEFINE_ECDSA(ECDSA_EC_PMUL_Solver, G1_1_G1SM2, Fq_SM2_1, Fq_SM2_n); +DEFINE_SECP256K1_FP(Fq_SECP256K1_1, FqSECP256K1, u32, 32, LayoutT<1>, 8, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SECP256K1); +DEFINE_FP(Fq_SECP256K1_n, FqSECP256K1_n, u32, 32, LayoutT<1>, 8); +DEFINE_EC(G1_1, G1SECP256K1, Fq_SECP256K1_1, SECP256K1_CURVE, 2); +DEFINE_ECDSA(ECDSA_EC_PMUL_Solver, G1_1_G1SECP256K1, Fq_SECP256K1_1, Fq_SECP256K1_n); TEST(ECDSA_EC_PMUL, Performance) { test_modinv_in_data_parallel(); } \ No newline at end of file