From b7240b3d8daf649f5dced716f0620cf07212fd92 Mon Sep 17 00:00:00 2001 From: rcxpony Date: Sun, 16 Mar 2025 22:03:05 +0500 Subject: [PATCH] optimized --- libs/fe.cu | 1187 ++++++++++++++++++----------------------------- libs/fe.cuh | 31 +- libs/fixedint.h | 51 -- libs/ge.cu | 270 ++--------- libs/ge.cuh | 15 - sources/main.cu | 33 +- 6 files changed, 510 insertions(+), 1077 deletions(-) delete mode 100644 libs/fixedint.h diff --git a/libs/fe.cu b/libs/fe.cu index bd7ff06..6b6593a 100644 --- a/libs/fe.cu +++ b/libs/fe.cu @@ -1,75 +1,23 @@ -#include #include -static uint64_t __host__ __device__ load_3(const unsigned char* in) { - return (uint64_t)in[0] | - ((uint64_t)in[1] << 8) | - ((uint64_t)in[2] << 16); -} -static uint64_t __host__ __device__ load_4(const unsigned char* in) { - return (uint64_t)in[0] | - ((uint64_t)in[1] << 8) | - ((uint64_t)in[2] << 16) | - ((uint64_t)in[3] << 24); -} -void __host__ __device__ fe_0(fe h) { -#pragma unroll 10 - for (int i = 0; i < 10; i++) h[i] = 0; -} void __host__ __device__ fe_1(fe h) { h[0] = 1; #pragma unroll 10 for (int i = 1; i < 10; i++) h[i] = 0; } -void __host__ __device__ fe_add(fe h, const fe f, const fe g) { +void __host__ __device__ fe_add(int h[10], const fe& f, const fe& g) { #pragma unroll 10 for (int i = 0; i < 10; i++) h[i] = f[i] + g[i]; } -void __host__ __device__ fe_cmov(fe f, const fe g, unsigned int b) { - int32_t mask = -((int)b); +void __host__ __device__ fe_cmov(fe f, const fe& g, unsigned int b) { + int mask = -((int)b); +#pragma unroll 10 for (int i = 0; i < 10; i++) f[i] ^= mask & (f[i] ^ g[i]); } -void fe_cswap(fe f, fe g, unsigned int b) { - int i, x; - b = -(int)b; - for (i = 0; i < 10; i++) { - x = f[i] ^ g[i]; - x &= b; - f[i] ^= x; - g[i] ^= x; - } -} -void __host__ __device__ fe_copy(fe h, const fe f) { +void __host__ __device__ fe_copy(fe h, const fe& f) { #pragma unroll 10 for (int i = 0; i < 10; i++) h[i] = f[i]; } -void __device__ __host__ fe_frombytes(fe h, const unsigned char* s) { - int64_t h0 = load_4(s); - int64_t h1 = load_3(s + 4) << 6; - int64_t h2 = load_3(s + 7) << 5; - int64_t h3 = load_3(s + 10) << 3; - int64_t h4 = load_3(s + 13) << 2; - int64_t h5 = load_4(s + 16); - int64_t h6 = load_3(s + 20) << 7; - int64_t h7 = load_3(s + 23) << 5; - int64_t h8 = load_3(s + 26) << 4; - int64_t h9 = (load_3(s + 29) & 8388607) << 2; - int64_t c; - c = (h9 + (1LL << 24)) >> 25; h0 += c * 19; h9 -= c << 25; - c = (h1 + (1LL << 24)) >> 25; h2 += c; h1 -= c << 25; - c = (h3 + (1LL << 24)) >> 25; h4 += c; h3 -= c << 25; - c = (h5 + (1LL << 24)) >> 25; h6 += c; h5 -= c << 25; - c = (h7 + (1LL << 24)) >> 25; h8 += c; h7 -= c << 25; - c = (h0 + (1LL << 25)) >> 26; h1 += c; h0 -= c << 26; - c = (h2 + (1LL << 25)) >> 26; h3 += c; h2 -= c << 26; - c = (h4 + (1LL << 25)) >> 26; h5 += c; h4 -= c << 26; - c = (h6 + (1LL << 25)) >> 26; h7 += c; h6 -= c << 26; - c = (h8 + (1LL << 25)) >> 26; h9 += c; h8 -= c << 26; - h[0] = (int32_t)h0; h[1] = (int32_t)h1; h[2] = (int32_t)h2; - h[3] = (int32_t)h3; h[4] = (int32_t)h4; h[5] = (int32_t)h5; - h[6] = (int32_t)h6; h[7] = (int32_t)h7; h[8] = (int32_t)h8; - h[9] = (int32_t)h9; -} -void fe_invert(fe out, const fe z) { +void fe_invert(fe out, const fe& z) { fe t0; fe t1; fe t2; @@ -142,613 +90,402 @@ void fe_invert(fe out, const fe z) { } fe_mul(out, t1, t0); } -int __host__ __device__ fe_isnegative(const fe f) { +int __host__ __device__ fe_isnegative(const fe& f) { unsigned char s[32]; fe_tobytes(s, f); return s[0] & 1; } -int __device__ __host__ fe_isnonzero(const fe f) { - unsigned char s[32]; - unsigned char r; - fe_tobytes(s, f); - r = s[0]; - #define F(i) r |= s[i] - F(1); - F(2); - F(3); - F(4); - F(5); - F(6); - F(7); - F(8); - F(9); - F(10); - F(11); - F(12); - F(13); - F(14); - F(15); - F(16); - F(17); - F(18); - F(19); - F(20); - F(21); - F(22); - F(23); - F(24); - F(25); - F(26); - F(27); - F(28); - F(29); - F(30); - F(31); - #undef F - return r != 0; -} -__device__ __host__ void fe_mul(fe h, const fe f, const fe g) { - const int32_t f0 = f[0], f1 = f[1], f2 = f[2], f3 = f[3], f4 = f[4], f5 = f[5], f6 = f[6], f7 = f[7], f8 = f[8], f9 = f[9]; - const int32_t g0 = g[0], g1 = g[1], g2 = g[2], g3 = g[3], g4 = g[4], g5 = g[5], g6 = g[6], g7 = g[7], g8 = g[8], g9 = g[9]; - const int32_t g1_19 = 19 * g1, - g2_19 = 19 * g2, - g3_19 = 19 * g3, - g4_19 = 19 * g4, - g5_19 = 19 * g5, - g6_19 = 19 * g6, - g7_19 = 19 * g7, - g8_19 = 19 * g8, - g9_19 = 19 * g9; - - const int32_t f1_2 = 2 * f1, f3_2 = 2 * f3, f5_2 = 2 * f5, f7_2 = 2 * f7, f9_2 = 2 * f9; - const int64_t f0g0 = (int64_t)f0 * g0; - const int64_t f0g1 = (int64_t)f0 * g1; - const int64_t f0g2 = (int64_t)f0 * g2; - const int64_t f0g3 = (int64_t)f0 * g3; - const int64_t f0g4 = (int64_t)f0 * g4; - const int64_t f0g5 = (int64_t)f0 * g5; - const int64_t f0g6 = (int64_t)f0 * g6; - const int64_t f0g7 = (int64_t)f0 * g7; - const int64_t f0g8 = (int64_t)f0 * g8; - const int64_t f0g9 = (int64_t)f0 * g9; - const int64_t f1g0 = (int64_t)f1 * g0; - const int64_t f1g1_2 = (int64_t)f1_2 * g1; - const int64_t f1g2 = (int64_t)f1 * g2; - const int64_t f1g3_2 = (int64_t)f1_2 * g3; - const int64_t f1g4 = (int64_t)f1 * g4; - const int64_t f1g5_2 = (int64_t)f1_2 * g5; - const int64_t f1g6 = (int64_t)f1 * g6; - const int64_t f1g7_2 = (int64_t)f1_2 * g7; - const int64_t f1g8 = (int64_t)f1 * g8; - const int64_t f1g9_38 = (int64_t)f1_2 * g9_19; - const int64_t f2g0 = (int64_t)f2 * g0; - const int64_t f2g1 = (int64_t)f2 * g1; - const int64_t f2g2 = (int64_t)f2 * g2; - const int64_t f2g3 = (int64_t)f2 * g3; - const int64_t f2g4 = (int64_t)f2 * g4; - const int64_t f2g5 = (int64_t)f2 * g5; - const int64_t f2g6 = (int64_t)f2 * g6; - const int64_t f2g7 = (int64_t)f2 * g7; - const int64_t f2g8_19 = (int64_t)f2 * g8_19; - const int64_t f2g9_19 = (int64_t)f2 * g9_19; - const int64_t f3g0 = (int64_t)f3 * g0; - const int64_t f3g1_2 = (int64_t)f3_2 * g1; - const int64_t f3g2 = (int64_t)f3 * g2; - const int64_t f3g3_2 = (int64_t)f3_2 * g3; - const int64_t f3g4 = (int64_t)f3 * g4; - const int64_t f3g5_2 = (int64_t)f3_2 * g5; - const int64_t f3g6 = (int64_t)f3 * g6; - const int64_t f3g7_38 = (int64_t)f3_2 * g7_19; - const int64_t f3g8_19 = (int64_t)f3 * g8_19; - const int64_t f3g9_38 = (int64_t)f3_2 * g9_19; - const int64_t f4g0 = (int64_t)f4 * g0; - const int64_t f4g1 = (int64_t)f4 * g1; - const int64_t f4g2 = (int64_t)f4 * g2; - const int64_t f4g3 = (int64_t)f4 * g3; - const int64_t f4g4 = (int64_t)f4 * g4; - const int64_t f4g5 = (int64_t)f4 * g5; - const int64_t f4g6_19 = (int64_t)f4 * g6_19; - const int64_t f4g7_19 = (int64_t)f4 * g7_19; - const int64_t f4g8_19 = (int64_t)f4 * g8_19; - const int64_t f4g9_19 = (int64_t)f4 * g9_19; - const int64_t f5g0 = (int64_t)f5 * g0; - const int64_t f5g1_2 = (int64_t)f5_2 * g1; - const int64_t f5g2 = (int64_t)f5 * g2; - const int64_t f5g3_2 = (int64_t)f5_2 * g3; - const int64_t f5g4 = (int64_t)f5 * g4; - const int64_t f5g5_38 = (int64_t)f5_2 * g5_19; - const int64_t f5g6_19 = (int64_t)f5 * g6_19; - const int64_t f5g7_38 = (int64_t)f5_2 * g7_19; - const int64_t f5g8_19 = (int64_t)f5 * g8_19; - const int64_t f5g9_38 = (int64_t)f5_2 * g9_19; - const int64_t f6g0 = (int64_t)f6 * g0; - const int64_t f6g1 = (int64_t)f6 * g1; - const int64_t f6g2 = (int64_t)f6 * g2; - const int64_t f6g3 = (int64_t)f6 * g3; - const int64_t f6g4_19 = (int64_t)f6 * g4_19; - const int64_t f6g5_19 = (int64_t)f6 * g5_19; - const int64_t f6g6_19 = (int64_t)f6 * g6_19; - const int64_t f6g7_19 = (int64_t)f6 * g7_19; - const int64_t f6g8_19 = (int64_t)f6 * g8_19; - const int64_t f6g9_19 = (int64_t)f6 * g9_19; - const int64_t f7g0 = (int64_t)f7 * g0; - const int64_t f7g1_2 = (int64_t)f7_2 * g1; - const int64_t f7g2 = (int64_t)f7 * g2; - const int64_t f7g3_38 = (int64_t)f7_2 * g3_19; - const int64_t f7g4_19 = (int64_t)f7 * g4_19; - const int64_t f7g5_38 = (int64_t)f7_2 * g5_19; - const int64_t f7g6_19 = (int64_t)f7 * g6_19; - const int64_t f7g7_38 = (int64_t)f7_2 * g7_19; - const int64_t f7g8_19 = (int64_t)f7 * g8_19; - const int64_t f7g9_38 = (int64_t)f7_2 * g9_19; - const int64_t f8g0 = (int64_t)f8 * g0; - const int64_t f8g1 = (int64_t)f8 * g1; - const int64_t f8g2_19 = (int64_t)f8 * g2_19; - const int64_t f8g3_19 = (int64_t)f8 * g3_19; - const int64_t f8g4_19 = (int64_t)f8 * g4_19; - const int64_t f8g5_19 = (int64_t)f8 * g5_19; - const int64_t f8g6_19 = (int64_t)f8 * g6_19; - const int64_t f8g7_19 = (int64_t)f8 * g7_19; - const int64_t f8g8_19 = (int64_t)f8 * g8_19; - const int64_t f8g9_19 = (int64_t)f8 * g9_19; - const int64_t f9g0 = (int64_t)f9 * g0; - const int64_t f9g1_38 = (int64_t)f9_2 * g1_19; - const int64_t f9g2_19 = (int64_t)f9 * g2_19; - const int64_t f9g3_38 = (int64_t)f9_2 * g3_19; - const int64_t f9g4_19 = (int64_t)f9 * g4_19; - const int64_t f9g5_38 = (int64_t)f9_2 * g5_19; - const int64_t f9g6_19 = (int64_t)f9 * g6_19; - const int64_t f9g7_38 = (int64_t)f9_2 * g7_19; - const int64_t f9g8_19 = (int64_t)f9 * g8_19; - const int64_t f9g9_38 = (int64_t)f9_2 * g9_19; - int64_t h0_val = f0g0 + f1g9_38 + f2g8_19 + f3g7_38 + f4g6_19 + f5g5_38 + f6g4_19 + f7g3_38 + f8g2_19 + f9g1_38; - int64_t h1_val = f0g1 + f1g0 + f2g9_19 + f3g8_19 + f4g7_19 + f5g6_19 + f6g5_19 + f7g4_19 + f8g3_19 + f9g2_19; - int64_t h2_val = f0g2 + f1g1_2 + f2g0 + f3g9_38 + f4g8_19 + f5g7_38 + f6g6_19 + f7g5_38 + f8g4_19 + f9g3_38; - int64_t h3_val = f0g3 + f1g2 + f2g1 + f3g0 + f4g9_19 + f5g8_19 + f6g7_19 + f7g6_19 + f8g5_19 + f9g4_19; - int64_t h4_val = f0g4 + f1g3_2 + f2g2 + f3g1_2 + f4g0 + f5g9_38 + f6g8_19 + f7g7_38 + f8g6_19 + f9g5_38; - int64_t h5_val = f0g5 + f1g4 + f2g3 + f3g2 + f4g1 + f5g0 + f6g9_19 + f7g8_19 + f8g7_19 + f9g6_19; - int64_t h6_val = f0g6 + f1g5_2 + f2g4 + f3g3_2 + f4g2 + f5g1_2 + f6g0 + f7g9_38 + f8g8_19 + f9g7_38; - int64_t h7_val = f0g7 + f1g6 + f2g5 + f3g4 + f4g3 + f5g2 + f6g1 + f7g0 + f8g9_19 + f9g8_19; - int64_t h8_val = f0g8 + f1g7_2 + f2g6 + f3g5_2 + f4g4 + f5g3_2 + f6g2 + f7g1_2 + f8g0 + f9g9_38; - int64_t h9_val = f0g9 + f1g8 + f2g7 + f3g6 + f4g5 + f5g4 + f6g3 + f7g2 + f8g1 + f9g0; - int64_t carry; - carry = (h0_val + (1LL << 25)) >> 26; +__device__ __host__ void fe_mul(fe h, const fe& f, const fe& g) { + const long f0g0 = static_cast(f[0]) * static_cast(g[0]); + const long f0g1 = static_cast(f[0]) * static_cast(g[1]); + const long f0g2 = static_cast(f[0]) * static_cast(g[2]); + const long f0g3 = static_cast(f[0]) * static_cast(g[3]); + const long f0g4 = static_cast(f[0]) * static_cast(g[4]); + const long f0g5 = static_cast(f[0]) * static_cast(g[5]); + const long f0g6 = static_cast(f[0]) * static_cast(g[6]); + const long f0g7 = static_cast(f[0]) * static_cast(g[7]); + const long f0g8 = static_cast(f[0]) * static_cast(g[8]); + const long f0g9 = static_cast(f[0]) * static_cast(g[9]); + const long f1g0 = static_cast(f[1]) * static_cast(g[0]); + const long f1g1_2 = static_cast(2 * f[1]) * static_cast(g[1]); + const long f1g2 = static_cast(f[1]) * static_cast(g[2]); + const long f1g3_2 = static_cast(2 * f[1]) * static_cast(g[3]); + const long f1g4 = static_cast(f[1]) * static_cast(g[4]); + const long f1g5_2 = static_cast(2 * f[1]) * static_cast(g[5]); + const long f1g6 = static_cast(f[1]) * static_cast(g[6]); + const long f1g7_2 = static_cast(2 * f[1]) * static_cast(g[7]); + const long f1g8 = static_cast(f[1]) * static_cast(g[8]); + const long f1g9_38 = static_cast(2 * f[1]) * static_cast(19 * g[9]); + const long f2g0 = static_cast(f[2]) * static_cast(g[0]); + const long f2g1 = static_cast(f[2]) * static_cast(g[1]); + const long f2g2 = static_cast(f[2]) * static_cast(g[2]); + const long f2g3 = static_cast(f[2]) * static_cast(g[3]); + const long f2g4 = static_cast(f[2]) * static_cast(g[4]); + const long f2g5 = static_cast(f[2]) * static_cast(g[5]); + const long f2g6 = static_cast(f[2]) * static_cast(g[6]); + const long f2g7 = static_cast(f[2]) * static_cast(g[7]); + const long f2g8_19 = static_cast(f[2]) * static_cast(19 * g[8]); + const long f2g9_19 = static_cast(f[2]) * static_cast(19 * g[9]); + const long f3g0 = static_cast(f[3]) * static_cast(g[0]); + const long f3g1_2 = static_cast(2 * f[3]) * static_cast(g[1]); + const long f3g2 = static_cast(f[3]) * static_cast(g[2]); + const long f3g3_2 = static_cast(2 * f[3]) * static_cast(g[3]); + const long f3g4 = static_cast(f[3]) * static_cast(g[4]); + const long f3g5_2 = static_cast(2 * f[3]) * static_cast(g[5]); + const long f3g6 = static_cast(f[3]) * static_cast(g[6]); + const long f3g7_38 = static_cast(2 * f[3]) * static_cast(19 * g[7]); + const long f3g8_19 = static_cast(f[3]) * static_cast(19 * g[8]); + const long f3g9_38 = static_cast(2 * f[3]) * static_cast(19 * g[9]); + const long f4g0 = static_cast(f[4]) * static_cast(g[0]); + const long f4g1 = static_cast(f[4]) * static_cast(g[1]); + const long f4g2 = static_cast(f[4]) * static_cast(g[2]); + const long f4g3 = static_cast(f[4]) * static_cast(g[3]); + const long f4g4 = static_cast(f[4]) * static_cast(g[4]); + const long f4g5 = static_cast(f[4]) * static_cast(g[5]); + const long f4g6_19 = static_cast(f[4]) * static_cast(19 * g[6]); + const long f4g7_19 = static_cast(f[4]) * static_cast(19 * g[7]); + const long f4g8_19 = static_cast(f[4]) * static_cast(19 * g[8]); + const long f4g9_19 = static_cast(f[4]) * static_cast(19 * g[9]); + const long f5g0 = static_cast(f[5]) * static_cast(g[0]); + const long f5g1_2 = static_cast(2 * f[5]) * static_cast(g[1]); + const long f5g2 = static_cast(f[5]) * static_cast(g[2]); + const long f5g3_2 = static_cast(2 * f[5]) * static_cast(g[3]); + const long f5g4 = static_cast(f[5]) * static_cast(g[4]); + const long f5g5_38 = static_cast(2 * f[5]) * static_cast(19 * g[5]); + const long f5g6_19 = static_cast(f[5]) * static_cast(19 * g[6]); + const long f5g7_38 = static_cast(2 * f[5]) * static_cast(19 * g[7]); + const long f5g8_19 = static_cast(f[5]) * static_cast(19 * g[8]); + const long f5g9_38 = static_cast(2 * f[5]) * static_cast(19 * g[9]); + const long f6g0 = static_cast(f[6]) * static_cast(g[0]); + const long f6g1 = static_cast(f[6]) * static_cast(g[1]); + const long f6g2 = static_cast(f[6]) * static_cast(g[2]); + const long f6g3 = static_cast(f[6]) * static_cast(g[3]); + const long f6g4_19 = static_cast(f[6]) * static_cast(19 * g[4]); + const long f6g5_19 = static_cast(f[6]) * static_cast(19 * g[5]); + const long f6g6_19 = static_cast(f[6]) * static_cast(19 * g[6]); + const long f6g7_19 = static_cast(f[6]) * static_cast(19 * g[7]); + const long f6g8_19 = static_cast(f[6]) * static_cast(19 * g[8]); + const long f6g9_19 = static_cast(f[6]) * static_cast(19 * g[9]); + const long f7g0 = static_cast(f[7]) * static_cast(g[0]); + const long f7g1_2 = static_cast(2 * f[7]) * static_cast(g[1]); + const long f7g2 = static_cast(f[7]) * static_cast(g[2]); + const long f7g3_38 = static_cast(2 * f[7]) * static_cast(19 * g[3]); + const long f7g4_19 = static_cast(f[7]) * static_cast(19 * g[4]); + const long f7g5_38 = static_cast(2 * f[7]) * static_cast(19 * g[5]); + const long f7g6_19 = static_cast(f[7]) * static_cast(19 * g[6]); + const long f7g7_38 = static_cast(2 * f[7]) * static_cast(19 * g[7]); + const long f7g8_19 = static_cast(f[7]) * static_cast(19 * g[8]); + const long f7g9_38 = static_cast(2 * f[7]) * static_cast(19 * g[9]); + const long f8g0 = static_cast(f[8]) * static_cast(g[0]); + const long f8g1 = static_cast(f[8]) * static_cast(g[1]); + const long f8g2_19 = static_cast(f[8]) * static_cast(19 * g[2]); + const long f8g3_19 = static_cast(f[8]) * static_cast(19 * g[3]); + const long f8g4_19 = static_cast(f[8]) * static_cast(19 * g[4]); + const long f8g5_19 = static_cast(f[8]) * static_cast(19 * g[5]); + const long f8g6_19 = static_cast(f[8]) * static_cast(19 * g[6]); + const long f8g7_19 = static_cast(f[8]) * static_cast(19 * g[7]); + const long f8g8_19 = static_cast(f[8]) * static_cast(19 * g[8]); + const long f8g9_19 = static_cast(f[8]) * static_cast(19 * g[9]); + const long f9g0 = static_cast(f[9]) * static_cast(g[0]); + const long f9g1_38 = static_cast(2 * f[9]) * static_cast(19 * g[1]); + const long f9g2_19 = static_cast(f[9]) * static_cast(19 * g[2]); + const long f9g3_38 = static_cast(2 * f[9]) * static_cast(19 * g[3]); + const long f9g4_19 = static_cast(f[9]) * static_cast(19 * g[4]); + const long f9g5_38 = static_cast(2 * f[9]) * static_cast(19 * g[5]); + const long f9g6_19 = static_cast(f[9]) * static_cast(19 * g[6]); + const long f9g7_38 = static_cast(2 * f[9]) * static_cast(19 * g[7]); + const long f9g8_19 = static_cast(f[9]) * static_cast(19 * g[8]); + const long f9g9_38 = static_cast(2 * f[9]) * static_cast(19 * g[9]); + long h0_val = f0g0 + f1g9_38 + f2g8_19 + f3g7_38 + f4g6_19 + f5g5_38 + f6g4_19 + f7g3_38 + f8g2_19 + f9g1_38; + long h1_val = f0g1 + f1g0 + f2g9_19 + f3g8_19 + f4g7_19 + f5g6_19 + f6g5_19 + f7g4_19 + f8g3_19 + f9g2_19; + long h2_val = f0g2 + f1g1_2 + f2g0 + f3g9_38 + f4g8_19 + f5g7_38 + f6g6_19 + f7g5_38 + f8g4_19 + f9g3_38; + long h3_val = f0g3 + f1g2 + f2g1 + f3g0 + f4g9_19 + f5g8_19 + f6g7_19 + f7g6_19 + f8g5_19 + f9g4_19; + long h4_val = f0g4 + f1g3_2 + f2g2 + f3g1_2 + f4g0 + f5g9_38 + f6g8_19 + f7g7_38 + f8g6_19 + f9g5_38; + long h5_val = f0g5 + f1g4 + f2g3 + f3g2 + f4g1 + f5g0 + f6g9_19 + f7g8_19 + f8g7_19 + f9g6_19; + long h6_val = f0g6 + f1g5_2 + f2g4 + f3g3_2 + f4g2 + f5g1_2 + f6g0 + f7g9_38 + f8g8_19 + f9g7_38; + long h7_val = f0g7 + f1g6 + f2g5 + f3g4 + f4g3 + f5g2 + f6g1 + f7g0 + f8g9_19 + f9g8_19; + long h8_val = f0g8 + f1g7_2 + f2g6 + f3g5_2 + f4g4 + f5g3_2 + f6g2 + f7g1_2 + f8g0 + f9g9_38; + long h9_val = f0g9 + f1g8 + f2g7 + f3g6 + f4g5 + f5g4 + f6g3 + f7g2 + f8g1 + f9g0; + long carry; + carry = (h0_val + (static_cast(1LL) << 25)) >> 26; h1_val += carry; h0_val -= carry << 26; - carry = (h4_val + (1LL << 25)) >> 26; + carry = (h4_val + (static_cast(1LL) << 25)) >> 26; h5_val += carry; h4_val -= carry << 26; - carry = (h1_val + (1LL << 24)) >> 25; + carry = (h1_val + (static_cast(1LL) << 24)) >> 25; h2_val += carry; h1_val -= carry << 25; - carry = (h5_val + (1LL << 24)) >> 25; + carry = (h5_val + (static_cast(1LL) << 24)) >> 25; h6_val += carry; h5_val -= carry << 25; - carry = (h2_val + (1LL << 25)) >> 26; + carry = (h2_val + (static_cast(1LL) << 25)) >> 26; h3_val += carry; h2_val -= carry << 26; - carry = (h6_val + (1LL << 25)) >> 26; + carry = (h6_val + (static_cast(1LL) << 25)) >> 26; h7_val += carry; h6_val -= carry << 26; - carry = (h3_val + (1LL << 24)) >> 25; + carry = (h3_val + (static_cast(1LL) << 24)) >> 25; h4_val += carry; h3_val -= carry << 25; - carry = (h7_val + (1LL << 24)) >> 25; + carry = (h7_val + (static_cast(1LL) << 24)) >> 25; h8_val += carry; h7_val -= carry << 25; - carry = (h4_val + (1LL << 25)) >> 26; + carry = (h4_val + (static_cast(1LL) << 25)) >> 26; h5_val += carry; h4_val -= carry << 26; - carry = (h8_val + (1LL << 25)) >> 26; + carry = (h8_val + (static_cast(1LL) << 25)) >> 26; h9_val += carry; h8_val -= carry << 26; - carry = (h9_val + (1LL << 24)) >> 25; + carry = (h9_val + (static_cast(1LL) << 24)) >> 25; h0_val += carry * 19; h9_val -= carry << 25; - carry = (h0_val + (1LL << 25)) >> 26; + carry = (h0_val + (static_cast(1LL) << 25)) >> 26; h1_val += carry; h0_val -= carry << 26; - h[0] = (int32_t)h0_val; - h[1] = (int32_t)h1_val; - h[2] = (int32_t)h2_val; - h[3] = (int32_t)h3_val; - h[4] = (int32_t)h4_val; - h[5] = (int32_t)h5_val; - h[6] = (int32_t)h6_val; - h[7] = (int32_t)h7_val; - h[8] = (int32_t)h8_val; - h[9] = (int32_t)h9_val; + h[0] = static_cast(h0_val); + h[1] = static_cast(h1_val); + h[2] = static_cast(h2_val); + h[3] = static_cast(h3_val); + h[4] = static_cast(h4_val); + h[5] = static_cast(h5_val); + h[6] = static_cast(h6_val); + h[7] = static_cast(h7_val); + h[8] = static_cast(h8_val); + h[9] = static_cast(h9_val); } -void fe_mul121666(fe h, fe f) { - int32_t f0 = f[0]; - int32_t f1 = f[1]; - int32_t f2 = f[2]; - int32_t f3 = f[3]; - int32_t f4 = f[4]; - int32_t f5 = f[5]; - int32_t f6 = f[6]; - int32_t f7 = f[7]; - int32_t f8 = f[8]; - int32_t f9 = f[9]; - int64_t h0 = f0 * (int64_t) 121666; - int64_t h1 = f1 * (int64_t) 121666; - int64_t h2 = f2 * (int64_t) 121666; - int64_t h3 = f3 * (int64_t) 121666; - int64_t h4 = f4 * (int64_t) 121666; - int64_t h5 = f5 * (int64_t) 121666; - int64_t h6 = f6 * (int64_t) 121666; - int64_t h7 = f7 * (int64_t) 121666; - int64_t h8 = f8 * (int64_t) 121666; - int64_t h9 = f9 * (int64_t) 121666; - int64_t carry0; - int64_t carry1; - int64_t carry2; - int64_t carry3; - int64_t carry4; - int64_t carry5; - int64_t carry6; - int64_t carry7; - int64_t carry8; - int64_t carry9; - carry9 = (h9 + (int64_t) (1<<24)) >> 25; h0 += carry9 * 19; h9 -= carry9 << 25; - carry1 = (h1 + (int64_t) (1<<24)) >> 25; h2 += carry1; h1 -= carry1 << 25; - carry3 = (h3 + (int64_t) (1<<24)) >> 25; h4 += carry3; h3 -= carry3 << 25; - carry5 = (h5 + (int64_t) (1<<24)) >> 25; h6 += carry5; h5 -= carry5 << 25; - carry7 = (h7 + (int64_t) (1<<24)) >> 25; h8 += carry7; h7 -= carry7 << 25; - carry0 = (h0 + (int64_t) (1<<25)) >> 26; h1 += carry0; h0 -= carry0 << 26; - carry2 = (h2 + (int64_t) (1<<25)) >> 26; h3 += carry2; h2 -= carry2 << 26; - carry4 = (h4 + (int64_t) (1<<25)) >> 26; h5 += carry4; h4 -= carry4 << 26; - carry6 = (h6 + (int64_t) (1<<25)) >> 26; h7 += carry6; h6 -= carry6 << 26; - carry8 = (h8 + (int64_t) (1<<25)) >> 26; h9 += carry8; h8 -= carry8 << 26; - h[0] = (int32_t) h0; - h[1] = (int32_t) h1; - h[2] = (int32_t) h2; - h[3] = (int32_t) h3; - h[4] = (int32_t) h4; - h[5] = (int32_t) h5; - h[6] = (int32_t) h6; - h[7] = (int32_t) h7; - h[8] = (int32_t) h8; - h[9] = (int32_t) h9; +void __device__ __host__ fe_neg(fe h, const fe& f) { + h[0] = -f[0]; + h[1] = -f[1]; + h[2] = -f[2]; + h[3] = -f[3]; + h[4] = -f[4]; + h[5] = -f[5]; + h[6] = -f[6]; + h[7] = -f[7]; + h[8] = -f[8]; + h[9] = -f[9]; } -void __device__ __host__ fe_neg(fe h, const fe f) { - int32_t f0 = f[0]; - int32_t f1 = f[1]; - int32_t f2 = f[2]; - int32_t f3 = f[3]; - int32_t f4 = f[4]; - int32_t f5 = f[5]; - int32_t f6 = f[6]; - int32_t f7 = f[7]; - int32_t f8 = f[8]; - int32_t f9 = f[9]; - int32_t h0 = -f0; - int32_t h1 = -f1; - int32_t h2 = -f2; - int32_t h3 = -f3; - int32_t h4 = -f4; - int32_t h5 = -f5; - int32_t h6 = -f6; - int32_t h7 = -f7; - int32_t h8 = -f8; - int32_t h9 = -f9; - h[0] = h0; - h[1] = h1; - h[2] = h2; - h[3] = h3; - h[4] = h4; - h[5] = h5; - h[6] = h6; - h[7] = h7; - h[8] = h8; - h[9] = h9; -} -void __device__ __host__ fe_pow22523(fe out, const fe z) { - fe t0; - fe t1; - fe t2; - int i; - fe_sq(t0, z); -#pragma unroll - for (i = 1; i < 1; ++i) { - fe_sq(t0, t0); - } - fe_sq(t1, t0); -#pragma unroll - for (i = 1; i < 2; ++i) { - fe_sq(t1, t1); - } - fe_mul(t1, z, t1); - fe_mul(t0, t0, t1); - fe_sq(t0, t0); -#pragma unroll - for (i = 1; i < 1; ++i) { - fe_sq(t0, t0); - } - fe_mul(t0, t1, t0); - fe_sq(t1, t0); -#pragma unroll - for (i = 1; i < 5; ++i) { - fe_sq(t1, t1); - } - fe_mul(t0, t1, t0); - fe_sq(t1, t0); -#pragma unroll - for (i = 1; i < 10; ++i) { - fe_sq(t1, t1); - } - fe_mul(t1, t1, t0); - fe_sq(t2, t1); -#pragma unroll - for (i = 1; i < 20; ++i) { - fe_sq(t2, t2); - } - fe_mul(t1, t2, t1); - fe_sq(t1, t1); -#pragma unroll - for (i = 1; i < 10; ++i) { - fe_sq(t1, t1); - } - fe_mul(t0, t1, t0); - fe_sq(t1, t0); -#pragma unroll - for (i = 1; i < 50; ++i) { - fe_sq(t1, t1); - } - fe_mul(t1, t1, t0); - fe_sq(t2, t1); -#pragma unroll - for (i = 1; i < 100; ++i) { - fe_sq(t2, t2); - } - fe_mul(t1, t2, t1); - fe_sq(t1, t1); -#pragma unroll - for (i = 1; i < 50; ++i) { - fe_sq(t1, t1); - } - fe_mul(t0, t1, t0); - fe_sq(t0, t0); -#pragma unroll - for (i = 1; i < 2; ++i) { - fe_sq(t0, t0); - } - fe_mul(out, t0, z); - return; -} -void __device__ __host__ fe_sq(fe h, const fe f) { - int32_t f0 = f[0]; - int32_t f1 = f[1]; - int32_t f2 = f[2]; - int32_t f3 = f[3]; - int32_t f4 = f[4]; - int32_t f5 = f[5]; - int32_t f6 = f[6]; - int32_t f7 = f[7]; - int32_t f8 = f[8]; - int32_t f9 = f[9]; - int32_t f0_2 = 2 * f0; - int32_t f1_2 = 2 * f1; - int32_t f2_2 = 2 * f2; - int32_t f3_2 = 2 * f3; - int32_t f4_2 = 2 * f4; - int32_t f5_2 = 2 * f5; - int32_t f6_2 = 2 * f6; - int32_t f7_2 = 2 * f7; - int32_t f5_38 = 38 * f5; - int32_t f6_19 = 19 * f6; - int32_t f7_38 = 38 * f7; - int32_t f8_19 = 19 * f8; - int32_t f9_38 = 38 * f9; - int64_t f0f0 = f0 * (int64_t) f0; - int64_t f0f1_2 = f0_2 * (int64_t) f1; - int64_t f0f2_2 = f0_2 * (int64_t) f2; - int64_t f0f3_2 = f0_2 * (int64_t) f3; - int64_t f0f4_2 = f0_2 * (int64_t) f4; - int64_t f0f5_2 = f0_2 * (int64_t) f5; - int64_t f0f6_2 = f0_2 * (int64_t) f6; - int64_t f0f7_2 = f0_2 * (int64_t) f7; - int64_t f0f8_2 = f0_2 * (int64_t) f8; - int64_t f0f9_2 = f0_2 * (int64_t) f9; - int64_t f1f1_2 = f1_2 * (int64_t) f1; - int64_t f1f2_2 = f1_2 * (int64_t) f2; - int64_t f1f3_4 = f1_2 * (int64_t) f3_2; - int64_t f1f4_2 = f1_2 * (int64_t) f4; - int64_t f1f5_4 = f1_2 * (int64_t) f5_2; - int64_t f1f6_2 = f1_2 * (int64_t) f6; - int64_t f1f7_4 = f1_2 * (int64_t) f7_2; - int64_t f1f8_2 = f1_2 * (int64_t) f8; - int64_t f1f9_76 = f1_2 * (int64_t) f9_38; - int64_t f2f2 = f2 * (int64_t) f2; - int64_t f2f3_2 = f2_2 * (int64_t) f3; - int64_t f2f4_2 = f2_2 * (int64_t) f4; - int64_t f2f5_2 = f2_2 * (int64_t) f5; - int64_t f2f6_2 = f2_2 * (int64_t) f6; - int64_t f2f7_2 = f2_2 * (int64_t) f7; - int64_t f2f8_38 = f2_2 * (int64_t) f8_19; - int64_t f2f9_38 = f2 * (int64_t) f9_38; - int64_t f3f3_2 = f3_2 * (int64_t) f3; - int64_t f3f4_2 = f3_2 * (int64_t) f4; - int64_t f3f5_4 = f3_2 * (int64_t) f5_2; - int64_t f3f6_2 = f3_2 * (int64_t) f6; - int64_t f3f7_76 = f3_2 * (int64_t) f7_38; - int64_t f3f8_38 = f3_2 * (int64_t) f8_19; - int64_t f3f9_76 = f3_2 * (int64_t) f9_38; - int64_t f4f4 = f4 * (int64_t) f4; - int64_t f4f5_2 = f4_2 * (int64_t) f5; - int64_t f4f6_38 = f4_2 * (int64_t) f6_19; - int64_t f4f7_38 = f4 * (int64_t) f7_38; - int64_t f4f8_38 = f4_2 * (int64_t) f8_19; - int64_t f4f9_38 = f4 * (int64_t) f9_38; - int64_t f5f5_38 = f5 * (int64_t) f5_38; - int64_t f5f6_38 = f5_2 * (int64_t) f6_19; - int64_t f5f7_76 = f5_2 * (int64_t) f7_38; - int64_t f5f8_38 = f5_2 * (int64_t) f8_19; - int64_t f5f9_76 = f5_2 * (int64_t) f9_38; - int64_t f6f6_19 = f6 * (int64_t) f6_19; - int64_t f6f7_38 = f6 * (int64_t) f7_38; - int64_t f6f8_38 = f6_2 * (int64_t) f8_19; - int64_t f6f9_38 = f6 * (int64_t) f9_38; - int64_t f7f7_38 = f7 * (int64_t) f7_38; - int64_t f7f8_38 = f7_2 * (int64_t) f8_19; - int64_t f7f9_76 = f7_2 * (int64_t) f9_38; - int64_t f8f8_19 = f8 * (int64_t) f8_19; - int64_t f8f9_38 = f8 * (int64_t) f9_38; - int64_t f9f9_38 = f9 * (int64_t) f9_38; - int64_t h0 = f0f0 + f1f9_76 + f2f8_38 + f3f7_76 + f4f6_38 + f5f5_38; - int64_t h1 = f0f1_2 + f2f9_38 + f3f8_38 + f4f7_38 + f5f6_38; - int64_t h2 = f0f2_2 + f1f1_2 + f3f9_76 + f4f8_38 + f5f7_76 + f6f6_19; - int64_t h3 = f0f3_2 + f1f2_2 + f4f9_38 + f5f8_38 + f6f7_38; - int64_t h4 = f0f4_2 + f1f3_4 + f2f2 + f5f9_76 + f6f8_38 + f7f7_38; - int64_t h5 = f0f5_2 + f1f4_2 + f2f3_2 + f6f9_38 + f7f8_38; - int64_t h6 = f0f6_2 + f1f5_4 + f2f4_2 + f3f3_2 + f7f9_76 + f8f8_19; - int64_t h7 = f0f7_2 + f1f6_2 + f2f5_2 + f3f4_2 + f8f9_38; - int64_t h8 = f0f8_2 + f1f7_4 + f2f6_2 + f3f5_4 + f4f4 + f9f9_38; - int64_t h9 = f0f9_2 + f1f8_2 + f2f7_2 + f3f6_2 + f4f5_2; - int64_t carry0; - int64_t carry1; - int64_t carry2; - int64_t carry3; - int64_t carry4; - int64_t carry5; - int64_t carry6; - int64_t carry7; - int64_t carry8; - int64_t carry9; - carry0 = (h0 + (int64_t) (1 << 25)) >> 26; +void __device__ __host__ fe_sq(fe h, const fe& f) { + int f2 = f[2]; + int f3 = f[3]; + int f4 = f[4]; + int f5 = f[5]; + int f6 = f[6]; + int f7 = f[7]; + int f8 = f[8]; + int f9 = f[9]; + int f2_2 = 2 * f2; + int f3_2 = 2 * f3; + int f4_2 = 2 * f4; + int f5_2 = 2 * f5; + int f6_2 = 2 * f6; + int f7_2 = 2 * f7; + int f5_38 = 38 * f5; + int f6_19 = 19 * f6; + int f7_38 = 38 * f7; + int f8_19 = 19 * f8; + int f9_38 = 38 * f9; + long f0f0 = f[0] * static_cast(f[0]); + long f0f1_2 = 2 * f[0] *static_cast(f[1]); + long f0f2_2 = 2 * f[0] * static_cast(f[2]); + long f0f3_2 = 2 * f[0] * static_cast(f[3]); + long f0f4_2 = 2 * f[0] * static_cast(f[4]); + long f0f5_2 = 2 * f[0] * static_cast(f[5]); + long f0f6_2 = 2 * f[0] * static_cast(f[6]); + long f0f7_2 = 2 * f[0] * static_cast(f[7]); + long f0f8_2 = 2 * f[0] * static_cast(f[8]); + long f0f9_2 = 2 * f[0] * static_cast(f[9]); + long f1f1_2 = 2 * f[1] * static_cast(f[1]); + long f1f2_2 = 2 * f[1] * static_cast(f[2]); + long f1f3_4 = 2 * f[1] * static_cast(f[3] * 2); + long f1f4_2 = 2 * f[1] * static_cast(f[4]); + long f1f5_4 = 2 * f[1] * static_cast(f[5] * 2); + long f1f6_2 = 2 * f[1] * static_cast(f[6]); + long f1f7_4 = 2 * f[1] * static_cast(f[7] * 2); + long f1f8_2 = 2 * f[1] * static_cast(f[8]); + long f1f9_76 = 2 * f[1] * static_cast(f[9] * 38); + long f2f2 = f2 * static_cast(f2); + long f2f3_2 = f2_2 * static_cast(f3); + long f2f4_2 = f2_2 * static_cast(f4); + long f2f5_2 = f2_2 * static_cast(f5); + long f2f6_2 = f2_2 * static_cast(f6); + long f2f7_2 = f2_2 * static_cast(f7); + long f2f8_38 = f2_2 * static_cast(f8_19); + long f2f9_38 = f2 * static_cast(f9_38); + long f3f3_2 = f3_2 * static_cast(f3); + long f3f4_2 = f3_2 * static_cast(f4); + long f3f5_4 = f3_2 * static_cast(f5_2); + long f3f6_2 = f3_2 * static_cast(f6); + long f3f7_76 = f3_2 * static_cast(f7_38); + long f3f8_38 = f3_2 * static_cast(f8_19); + long f3f9_76 = f3_2 * static_cast(f9_38); + long f4f4 = f4 * static_cast(f4); + long f4f5_2 = f4_2 * static_cast(f5); + long f4f6_38 = f4_2 * static_cast(f6_19); + long f4f7_38 = f4 * static_cast(f7_38); + long f4f8_38 = f4_2 * static_cast(f8_19); + long f4f9_38 = f4 * static_cast(f9_38); + long f5f5_38 = f5 * static_cast(f5_38); + long f5f6_38 = f5_2 * static_cast(f6_19); + long f5f7_76 = f5_2 * static_cast(f7_38); + long f5f8_38 = f5_2 * static_cast(f8_19); + long f5f9_76 = f5_2 * static_cast(f9_38); + long f6f6_19 = f6 * static_cast(f6_19); + long f6f7_38 = f6 * static_cast(f7_38); + long f6f8_38 = f6_2 * static_cast(f8_19); + long f6f9_38 = f6 * static_cast(f9_38); + long f7f7_38 = f7 * static_cast(f7_38); + long f7f8_38 = f7_2 * static_cast(f8_19); + long f7f9_76 = f7_2 * static_cast(f9_38); + long f8f8_19 = f8 * static_cast(f8_19); + long f8f9_38 = f8 * static_cast(f9_38); + long f9f9_38 = f9 * static_cast(f9_38); + long h0 = f0f0 + f1f9_76 + f2f8_38 + f3f7_76 + f4f6_38 + f5f5_38; + long h1 = f0f1_2 + f2f9_38 + f3f8_38 + f4f7_38 + f5f6_38; + long h2 = f0f2_2 + f1f1_2 + f3f9_76 + f4f8_38 + f5f7_76 + f6f6_19; + long h3 = f0f3_2 + f1f2_2 + f4f9_38 + f5f8_38 + f6f7_38; + long h4 = f0f4_2 + f1f3_4 + f2f2 + f5f9_76 + f6f8_38 + f7f7_38; + long h5 = f0f5_2 + f1f4_2 + f2f3_2 + f6f9_38 + f7f8_38; + long h6 = f0f6_2 + f1f5_4 + f2f4_2 + f3f3_2 + f7f9_76 + f8f8_19; + long h7 = f0f7_2 + f1f6_2 + f2f5_2 + f3f4_2 + f8f9_38; + long h8 = f0f8_2 + f1f7_4 + f2f6_2 + f3f5_4 + f4f4 + f9f9_38; + long h9 = f0f9_2 + f1f8_2 + f2f7_2 + f3f6_2 + f4f5_2; + long carry0, carry1, carry2, carry3, carry4, carry5, carry6, carry7, carry8, carry9; + carry0 = (h0 + static_cast(1 << 25)) >> 26; h1 += carry0; h0 -= carry0 << 26; - carry4 = (h4 + (int64_t) (1 << 25)) >> 26; + carry4 = (h4 + static_cast(1 << 25)) >> 26; h5 += carry4; h4 -= carry4 << 26; - carry1 = (h1 + (int64_t) (1 << 24)) >> 25; + carry1 = (h1 + static_cast(1 << 24)) >> 25; h2 += carry1; h1 -= carry1 << 25; - carry5 = (h5 + (int64_t) (1 << 24)) >> 25; + carry5 = (h5 + static_cast(1 << 24)) >> 25; h6 += carry5; h5 -= carry5 << 25; - carry2 = (h2 + (int64_t) (1 << 25)) >> 26; + carry2 = (h2 + static_cast(1 << 25)) >> 26; h3 += carry2; h2 -= carry2 << 26; - carry6 = (h6 + (int64_t) (1 << 25)) >> 26; + carry6 = (h6 + static_cast(1 << 25)) >> 26; h7 += carry6; h6 -= carry6 << 26; - carry3 = (h3 + (int64_t) (1 << 24)) >> 25; + carry3 = (h3 + static_cast(1 << 24)) >> 25; h4 += carry3; h3 -= carry3 << 25; - carry7 = (h7 + (int64_t) (1 << 24)) >> 25; + carry7 = (h7 + static_cast(1 << 24)) >> 25; h8 += carry7; h7 -= carry7 << 25; - carry4 = (h4 + (int64_t) (1 << 25)) >> 26; + carry4 = (h4 + static_cast(1 << 25)) >> 26; h5 += carry4; h4 -= carry4 << 26; - carry8 = (h8 + (int64_t) (1 << 25)) >> 26; + carry8 = (h8 + static_cast(1 << 25)) >> 26; h9 += carry8; h8 -= carry8 << 26; - carry9 = (h9 + (int64_t) (1 << 24)) >> 25; + carry9 = (h9 + static_cast(1 << 24)) >> 25; h0 += carry9 * 19; h9 -= carry9 << 25; - carry0 = (h0 + (int64_t) (1 << 25)) >> 26; + carry0 = (h0 + static_cast(1 << 25)) >> 26; h1 += carry0; h0 -= carry0 << 26; - h[0] = (int32_t) h0; - h[1] = (int32_t) h1; - h[2] = (int32_t) h2; - h[3] = (int32_t) h3; - h[4] = (int32_t) h4; - h[5] = (int32_t) h5; - h[6] = (int32_t) h6; - h[7] = (int32_t) h7; - h[8] = (int32_t) h8; - h[9] = (int32_t) h9; + h[0] = (int)h0; + h[1] = (int)h1; + h[2] = (int)h2; + h[3] = (int)h3; + h[4] = (int)h4; + h[5] = (int)h5; + h[6] = (int)h6; + h[7] = (int)h7; + h[8] = (int)h8; + h[9] = (int)h9; } -void __host__ __device__ fe_sq2(fe h, const fe f) { - int32_t f0 = f[0]; - int32_t f1 = f[1]; - int32_t f2 = f[2]; - int32_t f3 = f[3]; - int32_t f4 = f[4]; - int32_t f5 = f[5]; - int32_t f6 = f[6]; - int32_t f7 = f[7]; - int32_t f8 = f[8]; - int32_t f9 = f[9]; - int32_t f0_2 = 2 * f0; - int32_t f1_2 = 2 * f1; - int32_t f2_2 = 2 * f2; - int32_t f3_2 = 2 * f3; - int32_t f4_2 = 2 * f4; - int32_t f5_2 = 2 * f5; - int32_t f6_2 = 2 * f6; - int32_t f7_2 = 2 * f7; - int32_t f5_38 = 38 * f5; - int32_t f6_19 = 19 * f6; - int32_t f7_38 = 38 * f7; - int32_t f8_19 = 19 * f8; - int32_t f9_38 = 38 * f9; - int64_t f0f0 = f0 * (int64_t) f0; - int64_t f0f1_2 = f0_2 * (int64_t) f1; - int64_t f0f2_2 = f0_2 * (int64_t) f2; - int64_t f0f3_2 = f0_2 * (int64_t) f3; - int64_t f0f4_2 = f0_2 * (int64_t) f4; - int64_t f0f5_2 = f0_2 * (int64_t) f5; - int64_t f0f6_2 = f0_2 * (int64_t) f6; - int64_t f0f7_2 = f0_2 * (int64_t) f7; - int64_t f0f8_2 = f0_2 * (int64_t) f8; - int64_t f0f9_2 = f0_2 * (int64_t) f9; - int64_t f1f1_2 = f1_2 * (int64_t) f1; - int64_t f1f2_2 = f1_2 * (int64_t) f2; - int64_t f1f3_4 = f1_2 * (int64_t) f3_2; - int64_t f1f4_2 = f1_2 * (int64_t) f4; - int64_t f1f5_4 = f1_2 * (int64_t) f5_2; - int64_t f1f6_2 = f1_2 * (int64_t) f6; - int64_t f1f7_4 = f1_2 * (int64_t) f7_2; - int64_t f1f8_2 = f1_2 * (int64_t) f8; - int64_t f1f9_76 = f1_2 * (int64_t) f9_38; - int64_t f2f2 = f2 * (int64_t) f2; - int64_t f2f3_2 = f2_2 * (int64_t) f3; - int64_t f2f4_2 = f2_2 * (int64_t) f4; - int64_t f2f5_2 = f2_2 * (int64_t) f5; - int64_t f2f6_2 = f2_2 * (int64_t) f6; - int64_t f2f7_2 = f2_2 * (int64_t) f7; - int64_t f2f8_38 = f2_2 * (int64_t) f8_19; - int64_t f2f9_38 = f2 * (int64_t) f9_38; - int64_t f3f3_2 = f3_2 * (int64_t) f3; - int64_t f3f4_2 = f3_2 * (int64_t) f4; - int64_t f3f5_4 = f3_2 * (int64_t) f5_2; - int64_t f3f6_2 = f3_2 * (int64_t) f6; - int64_t f3f7_76 = f3_2 * (int64_t) f7_38; - int64_t f3f8_38 = f3_2 * (int64_t) f8_19; - int64_t f3f9_76 = f3_2 * (int64_t) f9_38; - int64_t f4f4 = f4 * (int64_t) f4; - int64_t f4f5_2 = f4_2 * (int64_t) f5; - int64_t f4f6_38 = f4_2 * (int64_t) f6_19; - int64_t f4f7_38 = f4 * (int64_t) f7_38; - int64_t f4f8_38 = f4_2 * (int64_t) f8_19; - int64_t f4f9_38 = f4 * (int64_t) f9_38; - int64_t f5f5_38 = f5 * (int64_t) f5_38; - int64_t f5f6_38 = f5_2 * (int64_t) f6_19; - int64_t f5f7_76 = f5_2 * (int64_t) f7_38; - int64_t f5f8_38 = f5_2 * (int64_t) f8_19; - int64_t f5f9_76 = f5_2 * (int64_t) f9_38; - int64_t f6f6_19 = f6 * (int64_t) f6_19; - int64_t f6f7_38 = f6 * (int64_t) f7_38; - int64_t f6f8_38 = f6_2 * (int64_t) f8_19; - int64_t f6f9_38 = f6 * (int64_t) f9_38; - int64_t f7f7_38 = f7 * (int64_t) f7_38; - int64_t f7f8_38 = f7_2 * (int64_t) f8_19; - int64_t f7f9_76 = f7_2 * (int64_t) f9_38; - int64_t f8f8_19 = f8 * (int64_t) f8_19; - int64_t f8f9_38 = f8 * (int64_t) f9_38; - int64_t f9f9_38 = f9 * (int64_t) f9_38; - int64_t h0 = f0f0 + f1f9_76 + f2f8_38 + f3f7_76 + f4f6_38 + f5f5_38; - int64_t h1 = f0f1_2 + f2f9_38 + f3f8_38 + f4f7_38 + f5f6_38; - int64_t h2 = f0f2_2 + f1f1_2 + f3f9_76 + f4f8_38 + f5f7_76 + f6f6_19; - int64_t h3 = f0f3_2 + f1f2_2 + f4f9_38 + f5f8_38 + f6f7_38; - int64_t h4 = f0f4_2 + f1f3_4 + f2f2 + f5f9_76 + f6f8_38 + f7f7_38; - int64_t h5 = f0f5_2 + f1f4_2 + f2f3_2 + f6f9_38 + f7f8_38; - int64_t h6 = f0f6_2 + f1f5_4 + f2f4_2 + f3f3_2 + f7f9_76 + f8f8_19; - int64_t h7 = f0f7_2 + f1f6_2 + f2f5_2 + f3f4_2 + f8f9_38; - int64_t h8 = f0f8_2 + f1f7_4 + f2f6_2 + f3f5_4 + f4f4 + f9f9_38; - int64_t h9 = f0f9_2 + f1f8_2 + f2f7_2 + f3f6_2 + f4f5_2; - int64_t carry0; - int64_t carry1; - int64_t carry2; - int64_t carry3; - int64_t carry4; - int64_t carry5; - int64_t carry6; - int64_t carry7; - int64_t carry8; - int64_t carry9; +void __host__ __device__ fe_sq2(fe h, const fe& f) { + int f0 = f[0]; + int f1 = f[1]; + int f2 = f[2]; + int f3 = f[3]; + int f4 = f[4]; + int f5 = f[5]; + int f6 = f[6]; + int f7 = f[7]; + int f8 = f[8]; + int f9 = f[9]; + int f0_2 = 2 * f0; + int f1_2 = 2 * f1; + int f2_2 = 2 * f2; + int f3_2 = 2 * f3; + int f4_2 = 2 * f4; + int f5_2 = 2 * f5; + int f6_2 = 2 * f6; + int f7_2 = 2 * f7; + int f5_38 = 38 * f5; + int f6_19 = 19 * f6; + int f7_38 = 38 * f7; + int f8_19 = 19 * f8; + int f9_38 = 38 * f9; + long f0f0 = f[0] * static_cast(f[0]); + long f0f1_2 = f0_2 * static_cast(f1); + long f0f2_2 = f0_2 * static_cast(f2); + long f0f3_2 = f0_2 * static_cast(f3); + long f0f4_2 = f0_2 * static_cast(f4); + long f0f5_2 = f0_2 * static_cast(f5); + long f0f6_2 = f0_2 * static_cast(f6); + long f0f7_2 = f0_2 * static_cast(f7); + long f0f8_2 = f0_2 * static_cast(f8); + long f0f9_2 = f0_2 * static_cast(f9); + long f1f1_2 = f1_2 * static_cast(f1); + long f1f2_2 = f1_2 * static_cast(f2); + long f1f3_4 = f1_2 * static_cast(f3_2); + long f1f4_2 = f1_2 * static_cast(f4); + long f1f5_4 = f1_2 * static_cast(f5_2); + long f1f6_2 = f1_2 * static_cast(f6); + long f1f7_4 = f1_2 * static_cast(f7_2); + long f1f8_2 = f1_2 * static_cast(f8); + long f1f9_76 = f1_2 * static_cast(f9_38); + long f2f2 = f2 * static_cast(f2); + long f2f3_2 = f2_2 * static_cast(f3); + long f2f4_2 = f2_2 * static_cast(f4); + long f2f5_2 = f2_2 * static_cast(f5); + long f2f6_2 = f2_2 * static_cast(f6); + long f2f7_2 = f2_2 * static_cast(f7); + long f2f8_38 = f2_2 * static_cast(f8_19); + long f2f9_38 = f2 * static_cast(f9_38); + long f3f3_2 = f3_2 * static_cast(f3); + long f3f4_2 = f3_2 * static_cast(f4); + long f3f5_4 = f3_2 * static_cast(f5_2); + long f3f6_2 = f3_2 * static_cast(f6); + long f3f7_76 = f3_2 * static_cast(f7_38); + long f3f8_38 = f3_2 * static_cast(f8_19); + long f3f9_76 = f3_2 * static_cast(f9_38); + long f4f4 = f4 * static_cast(f4); + long f4f5_2 = f4_2 * static_cast(f5); + long f4f6_38 = f4_2 * static_cast(f6_19); + long f4f7_38 = f4 * static_cast(f7_38); + long f4f8_38 = f4_2 * static_cast(f8_19); + long f4f9_38 = f4 * static_cast(f9_38); + long f5f5_38 = f5 * static_cast(f5_38); + long f5f6_38 = f5_2 * static_cast(f6_19); + long f5f7_76 = f5_2 * static_cast(f7_38); + long f5f8_38 = f5_2 * static_cast(f8_19); + long f5f9_76 = f5_2 * static_cast(f9_38); + long f6f6_19 = f6 * static_cast(f6_19); + long f6f7_38 = f6 * static_cast(f7_38); + long f6f8_38 = f6_2 * static_cast(f8_19); + long f6f9_38 = f6 * static_cast(f9_38); + long f7f7_38 = f7 * static_cast(f7_38); + long f7f8_38 = f7_2 * static_cast(f8_19); + long f7f9_76 = f7_2 * static_cast(f9_38); + long f8f8_19 = f8 * static_cast(f8_19); + long f8f9_38 = f8 * static_cast(f9_38); + long f9f9_38 = f9 * static_cast(f9_38); + long h0 = f0f0 + f1f9_76 + f2f8_38 + f3f7_76 + f4f6_38 + f5f5_38; + long h1 = f0f1_2 + f2f9_38 + f3f8_38 + f4f7_38 + f5f6_38; + long h2 = f0f2_2 + f1f1_2 + f3f9_76 + f4f8_38 + f5f7_76 + f6f6_19; + long h3 = f0f3_2 + f1f2_2 + f4f9_38 + f5f8_38 + f6f7_38; + long h4 = f0f4_2 + f1f3_4 + f2f2 + f5f9_76 + f6f8_38 + f7f7_38; + long h5 = f0f5_2 + f1f4_2 + f2f3_2 + f6f9_38 + f7f8_38; + long h6 = f0f6_2 + f1f5_4 + f2f4_2 + f3f3_2 + f7f9_76 + f8f8_19; + long h7 = f0f7_2 + f1f6_2 + f2f5_2 + f3f4_2 + f8f9_38; + long h8 = f0f8_2 + f1f7_4 + f2f6_2 + f3f5_4 + f4f4 + f9f9_38; + long h9 = f0f9_2 + f1f8_2 + f2f7_2 + f3f6_2 + f4f5_2; + long carry0; + long carry1; + long carry2; + long carry3; + long carry4; + long carry5; + long carry6; + long carry7; + long carry8; + long carry9; h0 += h0; h1 += h1; h2 += h2; @@ -759,118 +496,88 @@ void __host__ __device__ fe_sq2(fe h, const fe f) { h7 += h7; h8 += h8; h9 += h9; - carry0 = (h0 + (int64_t) (1 << 25)) >> 26; + carry0 = (h0 + static_cast(1 << 25)) >> 26; h1 += carry0; h0 -= carry0 << 26; - carry4 = (h4 + (int64_t) (1 << 25)) >> 26; + carry4 = (h4 + static_cast(1 << 25)) >> 26; h5 += carry4; h4 -= carry4 << 26; - carry1 = (h1 + (int64_t) (1 << 24)) >> 25; + carry1 = (h1 + static_cast(1 << 24)) >> 25; h2 += carry1; h1 -= carry1 << 25; - carry5 = (h5 + (int64_t) (1 << 24)) >> 25; + carry5 = (h5 + static_cast(1 << 24)) >> 25; h6 += carry5; h5 -= carry5 << 25; - carry2 = (h2 + (int64_t) (1 << 25)) >> 26; + carry2 = (h2 + static_cast(1 << 25)) >> 26; h3 += carry2; h2 -= carry2 << 26; - carry6 = (h6 + (int64_t) (1 << 25)) >> 26; + carry6 = (h6 + static_cast(1 << 25)) >> 26; h7 += carry6; h6 -= carry6 << 26; - carry3 = (h3 + (int64_t) (1 << 24)) >> 25; + carry3 = (h3 + static_cast(1 << 24)) >> 25; h4 += carry3; h3 -= carry3 << 25; - carry7 = (h7 + (int64_t) (1 << 24)) >> 25; + carry7 = (h7 + static_cast(1 << 24)) >> 25; h8 += carry7; h7 -= carry7 << 25; - carry4 = (h4 + (int64_t) (1 << 25)) >> 26; + carry4 = (h4 + static_cast(1 << 25)) >> 26; h5 += carry4; h4 -= carry4 << 26; - carry8 = (h8 + (int64_t) (1 << 25)) >> 26; + carry8 = (h8 + static_cast(1 << 25)) >> 26; h9 += carry8; h8 -= carry8 << 26; - carry9 = (h9 + (int64_t) (1 << 24)) >> 25; + carry9 = (h9 + static_cast(1 << 24)) >> 25; h0 += carry9 * 19; h9 -= carry9 << 25; - carry0 = (h0 + (int64_t) (1 << 25)) >> 26; + carry0 = (h0 + static_cast(1 << 25)) >> 26; h1 += carry0; h0 -= carry0 << 26; - h[0] = (int32_t) h0; - h[1] = (int32_t) h1; - h[2] = (int32_t) h2; - h[3] = (int32_t) h3; - h[4] = (int32_t) h4; - h[5] = (int32_t) h5; - h[6] = (int32_t) h6; - h[7] = (int32_t) h7; - h[8] = (int32_t) h8; - h[9] = (int32_t) h9; + h[0] = (int)h0; + h[1] = (int)h1; + h[2] = (int)h2; + h[3] = (int)h3; + h[4] = (int)h4; + h[5] = (int)h5; + h[6] = (int)h6; + h[7] = (int)h7; + h[8] = (int)h8; + h[9] = (int)h9; } -void __device__ __host__ fe_sub(fe h, const fe f, const fe g) { - int32_t f0 = f[0]; - int32_t f1 = f[1]; - int32_t f2 = f[2]; - int32_t f3 = f[3]; - int32_t f4 = f[4]; - int32_t f5 = f[5]; - int32_t f6 = f[6]; - int32_t f7 = f[7]; - int32_t f8 = f[8]; - int32_t f9 = f[9]; - int32_t g0 = g[0]; - int32_t g1 = g[1]; - int32_t g2 = g[2]; - int32_t g3 = g[3]; - int32_t g4 = g[4]; - int32_t g5 = g[5]; - int32_t g6 = g[6]; - int32_t g7 = g[7]; - int32_t g8 = g[8]; - int32_t g9 = g[9]; - int32_t h0 = f0 - g0; - int32_t h1 = f1 - g1; - int32_t h2 = f2 - g2; - int32_t h3 = f3 - g3; - int32_t h4 = f4 - g4; - int32_t h5 = f5 - g5; - int32_t h6 = f6 - g6; - int32_t h7 = f7 - g7; - int32_t h8 = f8 - g8; - int32_t h9 = f9 - g9; - h[0] = h0; - h[1] = h1; - h[2] = h2; - h[3] = h3; - h[4] = h4; - h[5] = h5; - h[6] = h6; - h[7] = h7; - h[8] = h8; - h[9] = h9; +void __device__ __host__ fe_sub(fe h, const fe& f, const fe& g) { + h[0] = (f[0] - g[0]); + h[1] = (f[1] - g[1]); + h[2] = (f[2] - g[2]); + h[3] = (f[3] - g[3]); + h[4] = (f[4] - g[4]); + h[5] = (f[5] - g[5]); + h[6] = (f[6] - g[6]); + h[7] = (f[7] - g[7]); + h[8] = (f[8] - g[8]); + h[9] = (f[9] - g[9]); } -void __host__ __device__ fe_tobytes(unsigned char *s, const fe h) { - int32_t h0 = h[0]; - int32_t h1 = h[1]; - int32_t h2 = h[2]; - int32_t h3 = h[3]; - int32_t h4 = h[4]; - int32_t h5 = h[5]; - int32_t h6 = h[6]; - int32_t h7 = h[7]; - int32_t h8 = h[8]; - int32_t h9 = h[9]; - int32_t q; - int32_t carry0; - int32_t carry1; - int32_t carry2; - int32_t carry3; - int32_t carry4; - int32_t carry5; - int32_t carry6; - int32_t carry7; - int32_t carry8; - int32_t carry9; - q = (19 * h9 + (((int32_t) 1) << 24)) >> 25; +void __host__ __device__ fe_tobytes(unsigned char* s, const fe& h) { + int h0 = h[0]; + int h1 = h[1]; + int h2 = h[2]; + int h3 = h[3]; + int h4 = h[4]; + int h5 = h[5]; + int h6 = h[6]; + int h7 = h[7]; + int h8 = h[8]; + int h9 = h[9]; + int q; + int carry0; + int carry1; + int carry2; + int carry3; + int carry4; + int carry5; + int carry6; + int carry7; + int carry8; + int carry9; + q = (19 * h9 + (((int)1) << 24)) >> 25; q = (h0 + q) >> 26; q = (h1 + q) >> 25; q = (h2 + q) >> 26; @@ -911,36 +618,36 @@ void __host__ __device__ fe_tobytes(unsigned char *s, const fe h) { h8 -= carry8 << 26; carry9 = h9 >> 25; h9 -= carry9 << 25; - s[0] = (unsigned char) (h0 >> 0); - s[1] = (unsigned char) (h0 >> 8); - s[2] = (unsigned char) (h0 >> 16); - s[3] = (unsigned char) ((h0 >> 24) | (h1 << 2)); - s[4] = (unsigned char) (h1 >> 6); - s[5] = (unsigned char) (h1 >> 14); - s[6] = (unsigned char) ((h1 >> 22) | (h2 << 3)); - s[7] = (unsigned char) (h2 >> 5); - s[8] = (unsigned char) (h2 >> 13); - s[9] = (unsigned char) ((h2 >> 21) | (h3 << 5)); - s[10] = (unsigned char) (h3 >> 3); - s[11] = (unsigned char) (h3 >> 11); - s[12] = (unsigned char) ((h3 >> 19) | (h4 << 6)); - s[13] = (unsigned char) (h4 >> 2); - s[14] = (unsigned char) (h4 >> 10); - s[15] = (unsigned char) (h4 >> 18); - s[16] = (unsigned char) (h5 >> 0); - s[17] = (unsigned char) (h5 >> 8); - s[18] = (unsigned char) (h5 >> 16); - s[19] = (unsigned char) ((h5 >> 24) | (h6 << 1)); - s[20] = (unsigned char) (h6 >> 7); - s[21] = (unsigned char) (h6 >> 15); - s[22] = (unsigned char) ((h6 >> 23) | (h7 << 3)); - s[23] = (unsigned char) (h7 >> 5); - s[24] = (unsigned char) (h7 >> 13); - s[25] = (unsigned char) ((h7 >> 21) | (h8 << 4)); - s[26] = (unsigned char) (h8 >> 4); - s[27] = (unsigned char) (h8 >> 12); - s[28] = (unsigned char) ((h8 >> 20) | (h9 << 6)); - s[29] = (unsigned char) (h9 >> 2); - s[30] = (unsigned char) (h9 >> 10); - s[31] = (unsigned char) (h9 >> 18); + s[0] = (unsigned char)(h0 >> 0); + s[1] = (unsigned char)(h0 >> 8); + s[2] = (unsigned char)(h0 >> 16); + s[3] = (unsigned char)((h0 >> 24) | (h1 << 2)); + s[4] = (unsigned char)(h1 >> 6); + s[5] = (unsigned char)(h1 >> 14); + s[6] = (unsigned char)((h1 >> 22) | (h2 << 3)); + s[7] = (unsigned char)(h2 >> 5); + s[8] = (unsigned char)(h2 >> 13); + s[9] = (unsigned char)((h2 >> 21) | (h3 << 5)); + s[10] = (unsigned char)(h3 >> 3); + s[11] = (unsigned char)(h3 >> 11); + s[12] = (unsigned char)((h3 >> 19) | (h4 << 6)); + s[13] = (unsigned char)(h4 >> 2); + s[14] = (unsigned char)(h4 >> 10); + s[15] = (unsigned char)(h4 >> 18); + s[16] = (unsigned char)(h5 >> 0); + s[17] = (unsigned char)(h5 >> 8); + s[18] = (unsigned char)(h5 >> 16); + s[19] = (unsigned char)((h5 >> 24) | (h6 << 1)); + s[20] = (unsigned char)(h6 >> 7); + s[21] = (unsigned char)(h6 >> 15); + s[22] = (unsigned char)((h6 >> 23) | (h7 << 3)); + s[23] = (unsigned char)(h7 >> 5); + s[24] = (unsigned char)(h7 >> 13); + s[25] = (unsigned char)((h7 >> 21) | (h8 << 4)); + s[26] = (unsigned char)(h8 >> 4); + s[27] = (unsigned char)(h8 >> 12); + s[28] = (unsigned char)((h8 >> 20) | (h9 << 6)); + s[29] = (unsigned char)(h9 >> 2); + s[30] = (unsigned char)(h9 >> 10); + s[31] = (unsigned char)(h9 >> 18); } \ No newline at end of file diff --git a/libs/fe.cuh b/libs/fe.cuh index 95d6008..b2d9ceb 100644 --- a/libs/fe.cuh +++ b/libs/fe.cuh @@ -1,23 +1,16 @@ #ifndef __FE_H #define __FE_H -#include -typedef int32_t fe[10]; -void __host__ __device__ fe_0(fe h); +using fe = int[10]; void __device__ __host__ fe_1(fe h); -void __device__ __host__ fe_frombytes(fe h, const unsigned char *s); -void __device__ __host__ fe_tobytes(unsigned char *s, const fe h); -void __host__ __device__ fe_copy(fe h, const fe f); -int __host__ __device__ fe_isnegative(const fe f); -int __device__ __host__ fe_isnonzero(const fe f); -void __host__ __device__ fe_cmov(fe f, const fe g, unsigned int b); -void fe_cswap(fe f, fe g, unsigned int b); -void __device__ __host__ fe_neg(fe h, const fe f); -void __device__ __host__ fe_add(fe h, const fe f, const fe g); -void __device__ __host__ fe_invert(fe out, const fe z); -void __device__ __host__ fe_sq(fe h, const fe f); -void __host__ __device__ fe_sq2(fe h, const fe f); -void __device__ __host__ fe_mul(fe h, const fe f, const fe g); -void fe_mul121666(fe h, fe f); -void __device__ __host__ fe_pow22523(fe out, const fe z); -void __device__ __host__ fe_sub(fe h, const fe f, const fe g); +void __device__ __host__ fe_tobytes(unsigned char *s, const fe& h); +void __host__ __device__ fe_copy(fe h, const fe& f); +int __host__ __device__ fe_isnegative(const fe& f); +void __host__ __device__ fe_cmov(fe f, const fe& g, unsigned int b); +void __device__ __host__ fe_neg(fe h, const fe& f); +void __device__ __host__ fe_add(fe h, const fe& f, const fe& g); +void __device__ __host__ fe_invert(fe out, const fe& z); +void __device__ __host__ fe_sq(fe h, const fe& f); +void __host__ __device__ fe_sq2(fe h, const fe& f); +void __device__ __host__ fe_mul(fe h, const fe& f, const fe& g); +void __device__ __host__ fe_sub(fe h, const fe& f, const fe& g); #endif diff --git a/libs/fixedint.h b/libs/fixedint.h deleted file mode 100644 index 513f1b2..0000000 --- a/libs/fixedint.h +++ /dev/null @@ -1,51 +0,0 @@ -#if ((defined(__STDC__) && __STDC__ && __STDC_VERSION__ >= 199901L) || (defined(__WATCOMC__) && (defined(_STDINT_H_INCLUDED) || __WATCOMC__ >= 1250)) || (defined(__GNUC__) && (defined(_STDINT_H) || defined(_STDINT_H_) || defined(__UINT_FAST64_TYPE__)) )) && !defined(FIXEDINT_H_INCLUDED) - #include - #define FIXEDINT_H_INCLUDED - #if defined(__WATCOMC__) && __WATCOMC__ >= 1250 && !defined(UINT64_C) - #include - #define UINT64_C(x) (x + (UINT64_MAX - UINT64_MAX)) - #endif -#endif -#ifndef FIXEDINT_H_INCLUDED - #define FIXEDINT_H_INCLUDED - #include - #ifndef uint32_t - #if (ULONG_MAX == 0xffffffffUL) - typedef unsigned long uint32_t; - #elif (UINT_MAX == 0xffffffffUL) - typedef unsigned int uint32_t; - #elif (USHRT_MAX == 0xffffffffUL) - typedef unsigned short uint32_t; - #endif - #endif - #ifndef int32_t - #if (LONG_MAX == 0x7fffffffL) - typedef signed long int32_t; - #elif (INT_MAX == 0x7fffffffL) - typedef signed int int32_t; - #elif (SHRT_MAX == 0x7fffffffL) - typedef signed short int32_t; - #endif - #endif - #if (defined(__STDC__) && defined(__STDC_VERSION__) && __STDC__ && __STDC_VERSION__ >= 199901L) - typedef long long int64_t; - typedef unsigned long long uint64_t; - #define UINT64_C(v) v ##ULL - #define INT64_C(v) v ##LL - #elif defined(__GNUC__) - __extension__ typedef long long int64_t; - __extension__ typedef unsigned long long uint64_t; - #define UINT64_C(v) v ##ULL - #define INT64_C(v) v ##LL - #elif defined(__MWERKS__) || defined(__SUNPRO_C) || defined(__SUNPRO_CC) || defined(__APPLE_CC__) || defined(_LONG_LONG) || defined(_CRAYC) - typedef long long int64_t; - typedef unsigned long long uint64_t; - #define UINT64_C(v) v ##ULL - #define INT64_C(v) v ##LL - #elif (defined(__WATCOMC__) && defined(__WATCOM_INT64__)) || (defined(_MSC_VER) && _INTEGRAL_MAX_BITS >= 64) || (defined(__BORLANDC__) && __BORLANDC__ > 0x460) || defined(__alpha) || defined(__DECC) - typedef __int64 int64_t; - typedef unsigned __int64 uint64_t; - #define UINT64_C(v) v ##UI64 - #define INT64_C(v) v ##I64 - #endif -#endif diff --git a/libs/ge.cu b/libs/ge.cu index 29c9151..c7b3a98 100644 --- a/libs/ge.cu +++ b/libs/ge.cu @@ -1,158 +1,6 @@ #include #include -void __host__ __device__ ge_add(ge_p1p1 *r, const ge_p3 *p, const ge_cached *q) { - fe t0; - fe_add(r->X, p->Y, p->X); - fe_sub(r->Y, p->Y, p->X); - fe_mul(r->Z, r->X, q->YplusX); - fe_mul(r->Y, r->Y, q->YminusX); - fe_mul(r->T, q->T2d, p->T); - fe_mul(r->X, p->Z, q->Z); - fe_add(t0, r->X, r->X); - fe_sub(r->X, r->Z, r->Y); - fe_add(r->Y, r->Z, r->Y); - fe_add(r->Z, t0, r->T); - fe_sub(r->T, t0, r->T); -} -static void __host__ __device__ slide(signed char *r, const unsigned char *a) { - int i; - int b; - int k; -#pragma unroll 256 - for (i = 0; i < 256; ++i) { - r[i] = 1 & (a[i >> 3] >> (i & 7)); - } -#pragma unroll 256 - for (i = 0; i < 256; ++i) - if (r[i]) { - #pragma unroll - for (b = 1; b <= 6 && i + b < 256; ++b) { - if (r[i + b]) { - if (r[i] + (r[i + b] << b) <= 15) { - r[i] += r[i + b] << b; - r[i + b] = 0; - } else if (r[i] - (r[i + b] << b) >= -15) { - r[i] -= r[i + b] << b; - #pragma unroll - for (k = i + b; k < 256; ++k) { - if (!r[k]) { - r[k] = 1; - break; - } - - r[k] = 0; - } - } else { - break; - } - } - } - } -} -void __host__ __device__ ge_double_scalarmult_vartime(ge_p2 *r, const unsigned char *a, const ge_p3 *A, const unsigned char *b) { - signed char aslide[256]; - signed char bslide[256]; - ge_cached Ai[8]; - ge_p1p1 t; - ge_p3 u; - ge_p3 A2; - int i; - slide(aslide, a); - slide(bslide, b); - ge_p3_to_cached(&Ai[0], A); - ge_p3_dbl(&t, A); - ge_p1p1_to_p3(&A2, &t); - ge_add(&t, &A2, &Ai[0]); - ge_p1p1_to_p3(&u, &t); - ge_p3_to_cached(&Ai[1], &u); - ge_add(&t, &A2, &Ai[1]); - ge_p1p1_to_p3(&u, &t); - ge_p3_to_cached(&Ai[2], &u); - ge_add(&t, &A2, &Ai[2]); - ge_p1p1_to_p3(&u, &t); - ge_p3_to_cached(&Ai[3], &u); - ge_add(&t, &A2, &Ai[3]); - ge_p1p1_to_p3(&u, &t); - ge_p3_to_cached(&Ai[4], &u); - ge_add(&t, &A2, &Ai[4]); - ge_p1p1_to_p3(&u, &t); - ge_p3_to_cached(&Ai[5], &u); - ge_add(&t, &A2, &Ai[5]); - ge_p1p1_to_p3(&u, &t); - ge_p3_to_cached(&Ai[6], &u); - ge_add(&t, &A2, &Ai[6]); - ge_p1p1_to_p3(&u, &t); - ge_p3_to_cached(&Ai[7], &u); - ge_p2_0(r); - for (i = 255; i >= 0; --i) { - if (aslide[i] || bslide[i]) { - break; - } - } - for (; i >= 0; --i) { - ge_p2_dbl(&t, r); - if (aslide[i] > 0) { - ge_p1p1_to_p3(&u, &t); - ge_add(&t, &u, &Ai[aslide[i] / 2]); - } else if (aslide[i] < 0) { - ge_p1p1_to_p3(&u, &t); - ge_sub(&t, &u, &Ai[(-aslide[i]) / 2]); - } - - if (bslide[i] > 0) { - ge_p1p1_to_p3(&u, &t); - ge_madd(&t, &u, &Bi[bslide[i] / 2]); - } else if (bslide[i] < 0) { - ge_p1p1_to_p3(&u, &t); - ge_msub(&t, &u, &Bi[(-bslide[i]) / 2]); - } - - ge_p1p1_to_p2(r, &t); - } -} -static __constant__ __device__ fe d = { - -10913610, 13857413, -15372611, 6949391, 114729, -8787816, -6275908, -3247719, -18696448, -12055116 -}; -static __constant__ __device__ fe sqrtm1 = { - -32595792, -7943725, 9377950, 3500415, 12389472, -272473, -25146209, -2005654, 326686, 11406482 -}; -int __device__ __host__ ge_frombytes_negate_vartime(ge_p3 *h, const unsigned char *s) { - fe u; - fe v; - fe v3; - fe vxx; - fe check; - fe_frombytes(h->Y, s); - fe_1(h->Z); - fe_sq(u, h->Y); - fe_mul(v, u, d); - fe_sub(u, u, h->Z); - fe_add(v, v, h->Z); - fe_sq(v3, v); - fe_mul(v3, v3, v); - fe_sq(h->X, v3); - fe_mul(h->X, h->X, v); - fe_mul(h->X, h->X, u); - fe_pow22523(h->X, h->X); - fe_mul(h->X, h->X, v3); - fe_mul(h->X, h->X, u); - fe_sq(vxx, h->X); - fe_mul(vxx, vxx, v); - fe_sub(check, vxx, u); - if (fe_isnonzero(check)) { - fe_add(check, vxx, u); - if (fe_isnonzero(check)) { - return -1; - } - fe_mul(h->X, h->X, sqrtm1); - } - if (fe_isnegative(h->X) == (s[31] >> 7)) { - fe_neg(h->X, h->X); - } - fe_mul(h->T, h->X, h->Y); - return 0; -} -void __host__ __device__ ge_madd(ge_p1p1 *r, const ge_p3 *p, const ge_precomp *q) { +void __host__ __device__ ge_madd(ge_p1p1* r, const ge_p3* p, const ge_precomp* q) { fe t0; fe_add(r->X, p->Y, p->X); fe_sub(r->Y, p->Y, p->X); @@ -165,36 +13,18 @@ void __host__ __device__ ge_madd(ge_p1p1 *r, const ge_p3 *p, const ge_precomp *q fe_add(r->Z, t0, r->T); fe_sub(r->T, t0, r->T); } -void __host__ __device__ ge_msub(ge_p1p1 *r, const ge_p3 *p, const ge_precomp *q) { - fe t0; - fe_add(r->X, p->Y, p->X); - fe_sub(r->Y, p->Y, p->X); - fe_mul(r->Z, r->X, q->yminusx); - fe_mul(r->Y, r->Y, q->yplusx); - fe_mul(r->T, q->xy2d, p->T); - fe_add(t0, p->Z, p->Z); - fe_sub(r->X, r->Z, r->Y); - fe_add(r->Y, r->Z, r->Y); - fe_sub(r->Z, t0, r->T); - fe_add(r->T, t0, r->T); -} -void __host__ __device__ ge_p1p1_to_p2(ge_p2 *r, const ge_p1p1 *p) { +void __host__ __device__ ge_p1p1_to_p2(ge_p2* r, const ge_p1p1* p) { fe_mul(r->X, p->X, p->T); fe_mul(r->Y, p->Y, p->Z); fe_mul(r->Z, p->Z, p->T); } -void __host__ __device__ ge_p1p1_to_p3(ge_p3 *r, const ge_p1p1 *p) { +void __host__ __device__ ge_p1p1_to_p3(ge_p3* r, const ge_p1p1* p) { fe_mul(r->X, p->X, p->T); fe_mul(r->Y, p->Y, p->Z); fe_mul(r->Z, p->Z, p->T); fe_mul(r->T, p->X, p->Y); } -void __host__ __device__ ge_p2_0(ge_p2 *h) { - fe_0(h->X); - fe_1(h->Y); - fe_1(h->Z); -} -void __host__ __device__ ge_p2_dbl(ge_p1p1 *r, const ge_p2 *p) { +void __host__ __device__ ge_p2_dbl(ge_p1p1* r, const ge_p2* p) { fe t0; fe_sq(r->X, p->X); fe_sq(r->Z, p->Y); @@ -206,35 +36,15 @@ void __host__ __device__ ge_p2_dbl(ge_p1p1 *r, const ge_p2 *p) { fe_sub(r->X, t0, r->Y); fe_sub(r->T, r->T, r->Z); } -void __host__ __device__ ge_p3_0(ge_p3 *h) { - fe_0(h->X); - fe_1(h->Y); - fe_1(h->Z); - fe_0(h->T); -} -void __host__ __device__ ge_p3_dbl(ge_p1p1 *r, const ge_p3 *p) { +void __host__ __device__ ge_p3_dbl(ge_p1p1* r, const ge_p3* p) { ge_p2 q; - ge_p3_to_p2(&q, p); + fe_copy(q.X, p->X); + fe_copy(q.Y, p->Y); + fe_copy(q.Z, p->Z); ge_p2_dbl(r, &q); } -static __constant__ __device__ fe d2 = { - -21827239, -5839606, -30745221, 13898782, 229458, 15978800, -12551817, -6495438, 29715968, 9444199 -}; -void __host__ __device__ ge_p3_to_cached(ge_cached *r, const ge_p3 *p) { - fe_add(r->YplusX, p->Y, p->X); - fe_sub(r->YminusX, p->Y, p->X); - fe_copy(r->Z, p->Z); - fe_mul(r->T2d, p->T, d2); -} -void ge_p3_to_p2(ge_p2 *r, const ge_p3 *p) { - fe_copy(r->X, p->X); - fe_copy(r->Y, p->Y); - fe_copy(r->Z, p->Z); -} -void ge_p3_tobytes(unsigned char *s, const ge_p3 *h) { - fe recip; - fe x; - fe y; +void ge_p3_tobytes(unsigned char* s, const ge_p3* h) { + fe recip, x, y; fe_invert(recip, h->Z); fe_mul(x, h->X, recip); fe_mul(y, h->Y, recip); @@ -242,31 +52,26 @@ void ge_p3_tobytes(unsigned char *s, const ge_p3 *h) { s[31] ^= fe_isnegative(x) << 7; } static unsigned char __host__ __device__ equal(signed char b, signed char c) { - unsigned char ub = b; - unsigned char uc = c; - unsigned char x = ub ^ uc; - uint64_t y = x; - y -= 1; - y >>= 63; - return (unsigned char) y; -} -static unsigned char __host__ __device__ negative(signed char b) { - uint64_t x = b; + unsigned long x = b ^ c; + x -= 1; x >>= 63; - return (unsigned char) x; + return (unsigned char)x; } -static void __host__ __device__ cmov(ge_precomp *t, const ge_precomp *u, unsigned char b) { +static void __host__ __device__ cmov(ge_precomp* t, const ge_precomp* u, unsigned char b) { fe_cmov(t->yplusx, u->yplusx, b); fe_cmov(t->yminusx, u->yminusx, b); fe_cmov(t->xy2d, u->xy2d, b); } -static void __host__ __device__ select(ge_precomp *t, int pos, signed char b) { +static void __host__ __device__ select(ge_precomp* t, int pos, signed char b) { ge_precomp minust; - unsigned char bnegative = negative(b); + unsigned long x = b; + x >>= 63; + unsigned char bnegative = static_cast(x); unsigned char babs = b - (((-bnegative) & b) << 1); fe_1(t->yplusx); fe_1(t->yminusx); - fe_0(t->xy2d); +#pragma unroll 10 + for (int i = 0; i < 10; i++) t->xy2d[i] = 0; cmov(t, &base[pos][0], equal(babs, 1)); cmov(t, &base[pos][1], equal(babs, 2)); cmov(t, &base[pos][2], equal(babs, 3)); @@ -298,8 +103,13 @@ void __device__ __host__ ge_scalarmult_base(ge_p3* h, const unsigned char* a) { e[i] -= carry << 4; } e[63] += carry; - ge_p3_0(h); -#pragma unroll 32 +#pragma unroll 10 + for (int i = 0; i < 10; i++) h->X[i] = 0; + fe_1(h->Y); + fe_1(h->Z); +#pragma unroll 10 + for (int i = 0; i < 10; i++) h->T[i] = 0; +#pragma unroll 64 for (i = 1; i < 64; i += 2) { select(&t, i >> 1, e[i]); ge_madd(&r, h, &t); @@ -310,34 +120,10 @@ void __device__ __host__ ge_scalarmult_base(ge_p3* h, const unsigned char* a) { ge_p1p1_to_p2(&s, &r); ge_p2_dbl(&r, &s); ge_p1p1_to_p2(&s, &r); ge_p2_dbl(&r, &s); ge_p1p1_to_p3(h, &r); -#pragma unroll 32 +#pragma unroll 64 for (i = 0; i < 64; i += 2) { select(&t, i >> 1, e[i]); ge_madd(&r, h, &t); ge_p1p1_to_p3(h, &r); } -} -void __host__ __device__ ge_sub(ge_p1p1 *r, const ge_p3 *p, const ge_cached *q) { - fe t0; - fe_add(r->X, p->Y, p->X); - fe_sub(r->Y, p->Y, p->X); - fe_mul(r->Z, r->X, q->YminusX); - fe_mul(r->Y, r->Y, q->YplusX); - fe_mul(r->T, q->T2d, p->T); - fe_mul(r->X, p->Z, q->Z); - fe_add(t0, r->X, r->X); - fe_sub(r->X, r->Z, r->Y); - fe_add(r->Y, r->Z, r->Y); - fe_sub(r->Z, t0, r->T); - fe_add(r->T, t0, r->T); -} -void __host__ __device__ ge_tobytes(unsigned char *s, const ge_p2 *h) { - fe recip; - fe x; - fe y; - fe_invert(recip, h->Z); - fe_mul(x, h->X, recip); - fe_mul(y, h->Y, recip); - fe_tobytes(s, y); - s[31] ^= fe_isnegative(x) << 7; } \ No newline at end of file diff --git a/libs/ge.cuh b/libs/ge.cuh index fb7c667..cd4abae 100644 --- a/libs/ge.cuh +++ b/libs/ge.cuh @@ -23,27 +23,12 @@ typedef struct { fe yminusx; fe xy2d; } ge_precomp; -typedef struct { - fe YplusX; - fe YminusX; - fe Z; - fe T2d; -} ge_cached; void __host__ __device__ ge_p3_tobytes(unsigned char *s, const ge_p3 *h); -void __host__ __device__ ge_tobytes(unsigned char *s, const ge_p2 *h); -int __host__ __device__ ge_frombytes_negate_vartime(ge_p3 *h, const unsigned char *s); -void __host__ __device__ ge_add(ge_p1p1 *r, const ge_p3 *p, const ge_cached *q); -void __host__ __device__ ge_sub(ge_p1p1 *r, const ge_p3 *p, const ge_cached *q); -void __host__ __device__ ge_double_scalarmult_vartime(ge_p2 *r, const unsigned char *a, const ge_p3 *A, const unsigned char *b); void __host__ __device__ ge_madd(ge_p1p1 *r, const ge_p3 *p, const ge_precomp *q); -void __host__ __device__ ge_msub(ge_p1p1 *r, const ge_p3 *p, const ge_precomp *q); void __host__ __device__ ge_scalarmult_base(ge_p3 *h, const unsigned char *a); void __host__ __device__ ge_p1p1_to_p2(ge_p2 *r, const ge_p1p1 *p); void __host__ __device__ ge_p1p1_to_p3(ge_p3 *r, const ge_p1p1 *p); -void __host__ __device__ ge_p2_0(ge_p2 *h); void __host__ __device__ ge_p2_dbl(ge_p1p1 *r, const ge_p2 *p); -void __host__ __device__ ge_p3_0(ge_p3 *h); void __host__ __device__ ge_p3_dbl(ge_p1p1 *r, const ge_p3 *p); -void __host__ __device__ ge_p3_to_cached(ge_cached *r, const ge_p3 *p); void __host__ __device__ ge_p3_to_p2(ge_p2 *r, const ge_p3 *p); #endif \ No newline at end of file diff --git a/sources/main.cu b/sources/main.cu index bfd04ec..8a4a34c 100644 --- a/sources/main.cu +++ b/sources/main.cu @@ -6,6 +6,13 @@ #include #include #include +#ifndef DEBUG + #define THREADSPB 256 + #define THDIVTHPB (tTh / THREADSPB) +#else + #define THREADSPB 1 + #define THDIVTHPB 1 +#endif __device__ unsigned d_high = 0x10; __device__ int parameters(const char* arg) noexcept { if ((cstring_find(arg, "--altitude") == 0 && cstring_length(arg) == 10) || @@ -114,9 +121,13 @@ __global__ void KeyGenKernel(curandState* randStates) { printf("\nIPv6:\t%s\nPK:\t%s\nSK:\t%s\n", getAddr(raw).data, ktos(keys.PublicKey).data, ktos(keys.PrivateKey).data); d_high = zeros; } + #ifdef DEBUG + if ((x & 0xFF) == 0) { + printf("Iters: %d\n", x); + } + #endif } } -#define THREADS_P_B 256 int main(int argc, char* argv[]) { int* d_result, mBpSM, h_high; char** d_argv; @@ -135,16 +146,17 @@ int main(int argc, char* argv[]) { cudaDeviceSynchronize(); cudaMemcpyFromSymbol(&h_high, d_high, sizeof(unsigned)); cudaGetDeviceProperties(&prop, 0); - cudaOccupancyMaxActiveBlocksPerMultiprocessor(&mBpSM, KeyGenKernel, THREADS_P_B, 0); - const int tTh = mBpSM * prop.multiProcessorCount * THREADS_P_B; - printf("High addrs: 2%02x+\nSMs: %d\nTotalThreads: %d\nBlocks: %d (Threads: %d)\n", h_high, prop.multiProcessorCount, tTh, tTh / THREADS_P_B, THREADS_P_B); + cudaOccupancyMaxActiveBlocksPerMultiprocessor(&mBpSM, KeyGenKernel, THREADSPB, 0); + const int tTh = mBpSM * prop.multiProcessorCount * THREADSPB; + printf("High addrs: 2%02x+\nSMs: %d\nTotalThreads: %d\nBlocks: %d (Threads: %d)\n", h_high, prop.multiProcessorCount, tTh, tTh / THREADSPB, THREADSPB); cudaMalloc(&rst, tTh * sizeof(curandState)); - unsigned int* d_seeds; - cudaMalloc(&d_seeds, tTh * sizeof(unsigned int)); - initRand<<>>(rst, d_seeds); + unsigned* d_seeds; + cudaMalloc(&d_seeds, tTh * sizeof(unsigned)); + initRand<<>>(rst, d_seeds); cudaDeviceSynchronize(); - unsigned int* h_seeds = (unsigned int*)malloc(tTh * sizeof(unsigned int)); - cudaMemcpy(h_seeds, d_seeds, tTh * sizeof(unsigned int), cudaMemcpyDeviceToHost); +#ifndef DEBUG + unsigned* h_seeds = (unsigned*)malloc(tTh * sizeof(unsigned)); + cudaMemcpy(h_seeds, d_seeds, tTh * sizeof(unsigned), cudaMemcpyDeviceToHost); if (checkSeeds(h_seeds, tTh)) { fprintf(stderr, "Duplicate seeds found!\n"); free(h_seeds); @@ -154,7 +166,8 @@ int main(int argc, char* argv[]) { } free(h_seeds); cudaFree(d_seeds); - KeyGenKernel << > > (rst); +#endif + KeyGenKernel<<>>(rst); cudaFree(rst); return 0; } \ No newline at end of file