diff --git a/.DS_Store b/.DS_Store deleted file mode 100644 index d11bd01..0000000 Binary files a/.DS_Store and /dev/null differ 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 ed7ee9a..0000000 Binary files a/include/.DS_Store and /dev/null differ diff --git a/include/gecc/.DS_Store b/include/gecc/.DS_Store deleted file mode 100644 index ca4c0ba..0000000 Binary files a/include/gecc/.DS_Store and /dev/null differ 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