diff --git a/libs/ed25519.cuh b/libs/ed25519.cuh index 2c791cb..a7ca1dc 100644 --- a/libs/ed25519.cuh +++ b/libs/ed25519.cuh @@ -48,27 +48,6 @@ __device__ __forceinline__ void ed25519_pack(unsigned char* c, const unsigned ch f25519_normalize(c); c[31] |= parity; } -/* -__device__ __forceinline__ unsigned char ed25519_try_unpack(unsigned char* x, unsigned char* y, const unsigned char* comp) { - int parity = comp[31] >> 7; - unsigned char a[F25519_SIZE], b[F25519_SIZE], c_[F25519_SIZE]; - f25519_copy(y, comp); - y[31] &= 127; - f25519_mul__distinct(c_, y, y); - f25519_mul__distinct(b, c_, ed25519_d); - f25519_add(a, b, f25519_one); - f25519_inv__distinct(b, a); - f25519_sub(a, c_, f25519_one); - f25519_mul__distinct(c_, a, b); - f25519_sqrt(a, c_); - f25519_neg(b, a); - f25519_select(x, a, b, (a[0] ^ parity) & 1); - f25519_mul__distinct(a, x, x); - f25519_normalize(a); - f25519_normalize(c_); - return f25519_eq(a, c_); -} -*/ __device__ __forceinline__ void ed25519_add(struct ed25519_pt* r, const struct ed25519_pt* p1, const struct ed25519_pt* p2) { unsigned char a[F25519_SIZE], b[F25519_SIZE], c[F25519_SIZE], d[F25519_SIZE]; unsigned char e[F25519_SIZE], f[F25519_SIZE], g[F25519_SIZE], h[F25519_SIZE]; diff --git a/libs/edsign.cuh b/libs/edsign.cuh index c20564a..0708fc1 100644 --- a/libs/edsign.cuh +++ b/libs/edsign.cuh @@ -2,12 +2,6 @@ #include #include #include -/* -__device__ __constant__ unsigned char ed25519_order[32] = { - 0xed, 0xd3, 0xf5, 0x5c, 0x1a, 0x63, 0x12, 0x58, 0xd6, 0x9c, 0xf7, 0xa2, 0xde, 0xf9, 0xde, 0x14, - 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x10 -}; -*/ __device__ __forceinline__ void expand_key(unsigned char* expanded, const unsigned char* secret) { struct sha512_state s; sha512_init(&s); @@ -15,14 +9,6 @@ __device__ __forceinline__ void expand_key(unsigned char* expanded, const unsign sha512_get(&s, expanded, 0, 64); ed25519_prepare(expanded); } -/* -__device__ __forceinline__ unsigned char upp(struct ed25519_pt* p, const unsigned char* packed) { - unsigned char x[F25519_SIZE], y[F25519_SIZE]; - unsigned char ok = ed25519_try_unpack(x, y, packed); - ed25519_project(p, x, y); - return ok; -} -*/ __device__ __forceinline__ void pp(unsigned char* packed, const struct ed25519_pt* p) { unsigned char x[F25519_SIZE], y[F25519_SIZE]; ed25519_unproject(x, y, p); @@ -37,60 +23,4 @@ __device__ __forceinline__ void edsign_sec_to_pub(unsigned char* pub, const unsi unsigned char expanded[64]; expand_key(expanded, secret); sm_pack(pub, expanded); -} -/* -__device__ __forceinline__ void hash_with_prefix(unsigned char* out_fp, unsigned char* init_block, unsigned int prefix_size, const unsigned char* message, unsigned long len) { - struct sha512_state s; - sha512_init(&s); - if (len < 128 && len + prefix_size < 128) { - memcpy(init_block + prefix_size, message, len); - sha512_final(&s, init_block, len + prefix_size); - } else { - unsigned long i; - memcpy(init_block + prefix_size, message, 128 - prefix_size); - sha512_block(&s, init_block); - for (i = 128 - prefix_size; i + 128 <= len; i += 128) { - sha512_block(&s, message + i); - } - sha512_final(&s, message + i, len - i + prefix_size); - } - sha512_get(&s, init_block, 0, 64); - fprime_from_bytes(out_fp, init_block, 64, ed25519_order); -} -__device__ __forceinline__ void generate_k(unsigned char* k, const unsigned char* kgen_key, const unsigned char* message, unsigned long len) { - unsigned char block[128]; - memcpy(block, kgen_key, 32); - hash_with_prefix(k, block, 32, message, len); -} -__device__ __forceinline__ void hash_message(unsigned char* z, const unsigned char* r, const unsigned char* a, const unsigned char* m, unsigned long len) { - unsigned char block[128]; - memcpy(block, r, 32); - memcpy(block + 32, a, 32); - hash_with_prefix(z, block, 64, m, len); -} -__device__ void edsign_sign(unsigned char* signature, const unsigned char* pub, const unsigned char* secret, const unsigned char* message, unsigned long len) { - unsigned char expanded[64]; - unsigned char e[32], s[32], k[32], z[32]; - expand_key(expanded, secret); - generate_k(k, expanded + 32, message, len); - sm_pack(signature, k); - hash_message(z, signature, pub, message, len); - fprime_from_bytes(e, expanded, 32, ed25519_order); - fprime_mul(s, z, e, ed25519_order); - fprime_add(s, k, ed25519_order); - memcpy(signature + 32, s, 32); -} -__device__ unsigned char edsign_verify(const unsigned char* signature, const unsigned char* pub, const unsigned char* message, unsigned long len) { - struct ed25519_pt p, q; - unsigned char lhs[F25519_SIZE], rhs[F25519_SIZE], z[32]; - unsigned char ok = 1; - hash_message(z, signature, pub, message, len); - sm_pack(lhs, signature + 32); - ok &= upp(&p, pub); - ed25519_smult(&p, &p, z); - ok &= upp(&q, signature); - ed25519_add(&p, &p, &q); - pp(rhs, &p); - return ok & f25519_eq(lhs, rhs); -} -*/ \ No newline at end of file +} \ No newline at end of file diff --git a/libs/f25519.cuh b/libs/f25519.cuh index 1635191..49e0a28 100644 --- a/libs/f25519.cuh +++ b/libs/f25519.cuh @@ -133,24 +133,6 @@ __device__ __forceinline__ void f25519_mul__distinct(unsigned char* __restrict__ c >>= 8; } } -/* -__device__ __forceinline__ void f25519_mul_c(unsigned char* __restrict__ r, const unsigned char* __restrict__ a, uint32_t b) { - uint32_t c = 0; -#pragma unroll - for (int i = 0; i < F25519_SIZE; i++) { - c = (c >> 8) + b * ((uint32_t)a[i]); - r[i] = (unsigned char)c; - } - r[F25519_SIZE - 1] &= 127; - c = (c >> 7) * 19; -#pragma unroll - for (int i = 0; i < F25519_SIZE; i++) { - c += r[i]; - r[i] = (unsigned char)c; - c >>= 8; - } -} -*/ __device__ __forceinline__ void f25519_inv__distinct(unsigned char* __restrict__ r, const unsigned char* __restrict__ x) { unsigned char s[F25519_SIZE]; f25519_mul__distinct(s, x, x); @@ -168,31 +150,4 @@ __device__ __forceinline__ void f25519_inv__distinct(unsigned char* __restrict__ f25519_mul__distinct(r, s, x); f25519_mul__distinct(s, r, r); f25519_mul__distinct(r, s, x); -} -/* -__device__ __forceinline__ void exp2523(unsigned char* __restrict__ r, const unsigned char* __restrict__ x, unsigned char* __restrict__ s) { - int i; - f25519_mul__distinct(r, x, x); - f25519_mul__distinct(s, r, x); -#pragma unroll - for (i = 0; i < 248; i++) { - f25519_mul__distinct(r, s, s); - f25519_mul__distinct(s, r, x); - } - f25519_mul__distinct(r, s, s); - f25519_mul__distinct(s, r, r); - f25519_mul__distinct(r, s, x); -} -__device__ __forceinline__ void f25519_sqrt(unsigned char* __restrict__ r, const unsigned char* __restrict__ a) { - unsigned char v[F25519_SIZE], i_val[F25519_SIZE], x[F25519_SIZE], y[F25519_SIZE]; - f25519_mul_c(x, a, 2); - exp2523(v, x, y); - f25519_mul__distinct(y, v, v); - f25519_mul__distinct(i_val, x, y); - unsigned char one[F25519_SIZE]; - f25519_load(one, 1); - f25519_sub(i_val, i_val, one); - f25519_mul__distinct(x, v, a); - f25519_mul__distinct(r, x, i_val); -} -*/ \ No newline at end of file +} \ No newline at end of file