sparrowwallet/gECC
sparrowwallet/gECC
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

adapt curve from sm2 to secp256k1

Browse Source
This commit is contained in:
Craig Raw 2025-10-15 13:44:34 +02:00
parent 322c1c143d
commit 200642b99f
18 changed files with 70 additions and 64 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

BIN
.DS_Store vendored
View File

Binary file not shown.

3
.gitignore vendored
View File

@ -17,4 +17,5 @@ __pycache__/
*$py.class
*.sage.py
.ipynb_checkpoints/
.ipynb_checkpoints/
.DS_Store

BIN
include/.DS_Store vendored
View File

Binary file not shown.

BIN
include/gecc/.DS_Store vendored
View File

Binary file not shown.

4
include/gecc/arith/details/fp_mont_multiply.h
View File

@ -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);

4
include/gecc/arith/fp.h
View File

@ -13,7 +13,7 @@
using FP_NAME = \
FpT<FP_NAME##Factory, gecc::arith::constants::FP_TYPE, FP_NAME##DCONST>
#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<gecc::arith::DigitT<DIGIT_TYPE, DIGIT_WIDTH>, 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 <typename D, typename L, u32 N> struct FpFactory {

19
scripts/constants.py
View File

@ -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,

16
scripts/constants_generator.py
View File

@ -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)

8
scripts/ec.py
View File

@ -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()

6
scripts/field.py
View File

@ -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)
Fq_SECP256K1 = Fp('FqSECP256K1', SECP256K1_q, 64, rexp=4)
Fq_SECP256K1_n = Fp('FqSECP256K1_n', SECP256K1_n, 64, rexp=4)

8
test/ecdsa_ec_fixed_pmul.cu
View File

@ -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<ECDSA_EC_PMUL_Solver>(); }

8
test/ecdsa_ec_unknown_pmul.cu
View File

@ -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<ECDSA_EC_PMUL_Solver>(); }

8
test/ecdsa_sign.cu
View File

@ -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<ECDSA_solver>(); }
TEST(ECDSA, Performance) { test_ecdsa_sign<ECDSA_solver>(); }

10
test/ecdsa_sign_baseline.cu
View File

@ -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<ECDSA_solver>(); }
TEST(ECDSA, Performance) { test_ecdsa_sign<ECDSA_solver>(); }

8
test/ecdsa_verify.cu
View File

@ -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<ECDSA_Verify_Solver>(); }
TEST(ECDSA, Performance) { test_ecdsa_verify<ECDSA_Verify_Solver>(); }

8
test/ecdsa_verify_baseline.cu
View File

@ -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<ECDSA_Verify_Solver>(); }
TEST(ECDSA, Performance) { test_ecdsa_verify<ECDSA_Verify_Solver>(); }

16
test/fp.cu
View File

@ -155,14 +155,14 @@ template <typename Field> 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<LAYOUT_WIDTH>, LIMBS, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SM2); \
TEST(FqSM2256K1_FP, FIELD##Correctness) { using namespace FqSM2_fp_test; test_fp<FIELD>(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<LAYOUT_WIDTH>, LIMBS, gecc::arith::MONTFLAG::SOS, gecc::arith::CURVEFLAG::SECP256K1); \
TEST(FqSECP256K1_FP, FIELD##Correctness) { using namespace FqSECP256K1_fp_test; test_fp<FIELD>(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<LAYOUT_WIDTH>, LIMBS); \
TEST(FqSM2_FP_n, FIELD##Correctness) { using namespace FqSM2_n_fp_test; test_fp<FIELD>(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<LAYOUT_WIDTH>, LIMBS); \
TEST(FqSECP256K1_FP_n, FIELD##Correctness) { using namespace FqSECP256K1_n_fp_test; test_fp<FIELD>(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)
ADD_FqSECP256K1_FP_TEST(Field_SECP256K1, u32, 32, 1, 8)
ADD_FqSECP256K1_n_FP_TEST(Field_SECP256K1_n, u32, 32, 1, 8)

8
test/modinv_data_parallel_profiling.cu
View File

@ -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<ECDSA_EC_PMUL_Solver>(); }