test

libs/ed25519.cu

@@ -83,7 +83,7 @@ __device__ __forceinline__ void ed25519_copy(struct ed25519_pt* __restrict__ dst
 }
 __device__ void ed25519_smult(ed25519_pt* r_out, const unsigned char* __restrict__ e) {
     ed25519_pt r = ed25519_neutral;
-#pragma unroll 256
+#pragma unroll
     for (int i = 255; i >= 0; i--) {
         struct ed25519_pt s;
         ed25519_double(&r, &r);

libs/edsign.cu

@@ -1,6 +1,7 @@
 #include <edsign.cuh>
 #include <ed25519.cuh>
 #include <sha512.cuh>
+#include <ge.cuh>
 __device__ __forceinline__ void expand_key(unsigned char* expanded, const unsigned char* secret) {
     struct sha512_state s;
     memcpy(&s, &sha512_initial_state, sizeof(s));
@@ -24,3 +25,12 @@ __device__ void ed25519_keygen(unsigned char private_key[64], unsigned char publ
     memcpy(private_key, random_seed, 32);
     memcpy(private_key + 32, public_key, 32);
 }
+__device__ void ed25519_create_keypair(unsigned char private_key[64], unsigned char public_key[32], unsigned char seed[32]) {
+    unsigned char expanded[64];
+    expand_key(expanded, seed);
+    ge_p3 A;
+    ge_scalarmult_base(&A, expanded);
+    ge_p3_tobytes(public_key, &A);
+    memcpy(private_key, seed, 32);
+    memcpy(private_key + 32, public_key, 32);
+}

libs/edsign.cuh

@@ -3,4 +3,5 @@
 __device__ __forceinline__ void expand_key(unsigned char* expanded, const unsigned char* secret);
 __device__ __forceinline__ void sm_pack(unsigned char* r, const unsigned char* k);
 __device__ void ed25519_keygen(unsigned char private_key[64], unsigned char public_key[32], unsigned char random_seed[32]);
+__device__ void ed25519_create_keypair(unsigned char private_key[64], unsigned char public_key[32], unsigned char seed[32]);
 #endif

libs/fe.cu

@@ -0,0 +1,946 @@
+#include <fixedint.h>
+#include <fe.cuh>
+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) {
+#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);
+    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) {
+#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) {
+    fe t0;
+    fe t1;
+    fe t2;
+    fe t3;
+    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(t2, t0);
+#pragma unroll
+    for (i = 1; i < 1; ++i) {
+        fe_sq(t2, t2);
+    }
+    fe_mul(t1, t1, t2);
+    fe_sq(t2, t1);
+#pragma unroll
+    for (i = 1; i < 5; ++i) {
+        fe_sq(t2, t2);
+    }
+    fe_mul(t1, t2, t1);
+    fe_sq(t2, t1);
+#pragma unroll
+    for (i = 1; i < 10; ++i) {
+        fe_sq(t2, t2);
+    }
+    fe_mul(t2, t2, t1);
+    fe_sq(t3, t2);
+#pragma unroll
+    for (i = 1; i < 20; ++i) {
+        fe_sq(t3, t3);
+    }
+    fe_mul(t2, t3, t2);
+    fe_sq(t2, t2);
+#pragma unroll
+    for (i = 1; i < 10; ++i) {
+        fe_sq(t2, t2);
+    }
+    fe_mul(t1, t2, t1);
+    fe_sq(t2, t1);
+#pragma unroll
+    for (i = 1; i < 50; ++i) {
+        fe_sq(t2, t2);
+    }
+    fe_mul(t2, t2, t1);
+    fe_sq(t3, t2);
+#pragma unroll
+    for (i = 1; i < 100; ++i) {
+        fe_sq(t3, t3);
+    }
+    fe_mul(t2, t3, t2);
+    fe_sq(t2, t2);
+#pragma unroll
+    for (i = 1; i < 50; ++i) {
+        fe_sq(t2, t2);
+    }
+    fe_mul(t1, t2, t1);
+    fe_sq(t1, t1);
+#pragma unroll
+    for (i = 1; i < 5; ++i) {
+        fe_sq(t1, t1);
+    }
+    fe_mul(out, t1, t0);
+}
+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;
+    h1_val += carry; h0_val -= carry << 26;
+    carry = (h4_val + (1LL << 25)) >> 26;
+    h5_val += carry; h4_val -= carry << 26;
+    carry = (h1_val + (1LL << 24)) >> 25;
+    h2_val += carry; h1_val -= carry << 25;
+    carry = (h5_val + (1LL << 24)) >> 25;
+    h6_val += carry; h5_val -= carry << 25;
+    carry = (h2_val + (1LL << 25)) >> 26;
+    h3_val += carry; h2_val -= carry << 26;
+    carry = (h6_val + (1LL << 25)) >> 26;
+    h7_val += carry; h6_val -= carry << 26;
+    carry = (h3_val + (1LL << 24)) >> 25;
+    h4_val += carry; h3_val -= carry << 25;
+    carry = (h7_val + (1LL << 24)) >> 25;
+    h8_val += carry; h7_val -= carry << 25;
+    carry = (h4_val + (1LL << 25)) >> 26;
+    h5_val += carry; h4_val -= carry << 26;
+    carry = (h8_val + (1LL << 25)) >> 26;
+    h9_val += carry; h8_val -= carry << 26;
+    carry = (h9_val + (1LL << 24)) >> 25;
+    h0_val += carry * 19; h9_val -= carry << 25;
+    carry = (h0_val + (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;
+}
+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) {
+    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;
+    h1 += carry0;
+    h0 -= carry0 << 26;
+    carry4 = (h4 + (int64_t) (1 << 25)) >> 26;
+    h5 += carry4;
+    h4 -= carry4 << 26;
+    carry1 = (h1 + (int64_t) (1 << 24)) >> 25;
+    h2 += carry1;
+    h1 -= carry1 << 25;
+    carry5 = (h5 + (int64_t) (1 << 24)) >> 25;
+    h6 += carry5;
+    h5 -= carry5 << 25;
+    carry2 = (h2 + (int64_t) (1 << 25)) >> 26;
+    h3 += carry2;
+    h2 -= carry2 << 26;
+    carry6 = (h6 + (int64_t) (1 << 25)) >> 26;
+    h7 += carry6;
+    h6 -= carry6 << 26;
+    carry3 = (h3 + (int64_t) (1 << 24)) >> 25;
+    h4 += carry3;
+    h3 -= carry3 << 25;
+    carry7 = (h7 + (int64_t) (1 << 24)) >> 25;
+    h8 += carry7;
+    h7 -= carry7 << 25;
+    carry4 = (h4 + (int64_t) (1 << 25)) >> 26;
+    h5 += carry4;
+    h4 -= carry4 << 26;
+    carry8 = (h8 + (int64_t) (1 << 25)) >> 26;
+    h9 += carry8;
+    h8 -= carry8 << 26;
+    carry9 = (h9 + (int64_t) (1 << 24)) >> 25;
+    h0 += carry9 * 19;
+    h9 -= carry9 << 25;
+    carry0 = (h0 + (int64_t) (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;
+}
+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;
+    h0 += h0;
+    h1 += h1;
+    h2 += h2;
+    h3 += h3;
+    h4 += h4;
+    h5 += h5;
+    h6 += h6;
+    h7 += h7;
+    h8 += h8;
+    h9 += h9;
+    carry0 = (h0 + (int64_t) (1 << 25)) >> 26;
+    h1 += carry0;
+    h0 -= carry0 << 26;
+    carry4 = (h4 + (int64_t) (1 << 25)) >> 26;
+    h5 += carry4;
+    h4 -= carry4 << 26;
+    carry1 = (h1 + (int64_t) (1 << 24)) >> 25;
+    h2 += carry1;
+    h1 -= carry1 << 25;
+    carry5 = (h5 + (int64_t) (1 << 24)) >> 25;
+    h6 += carry5;
+    h5 -= carry5 << 25;
+    carry2 = (h2 + (int64_t) (1 << 25)) >> 26;
+    h3 += carry2;
+    h2 -= carry2 << 26;
+    carry6 = (h6 + (int64_t) (1 << 25)) >> 26;
+    h7 += carry6;
+    h6 -= carry6 << 26;
+    carry3 = (h3 + (int64_t) (1 << 24)) >> 25;
+    h4 += carry3;
+    h3 -= carry3 << 25;
+    carry7 = (h7 + (int64_t) (1 << 24)) >> 25;
+    h8 += carry7;
+    h7 -= carry7 << 25;
+    carry4 = (h4 + (int64_t) (1 << 25)) >> 26;
+    h5 += carry4;
+    h4 -= carry4 << 26;
+    carry8 = (h8 + (int64_t) (1 << 25)) >> 26;
+    h9 += carry8;
+    h8 -= carry8 << 26;
+    carry9 = (h9 + (int64_t) (1 << 24)) >> 25;
+    h0 += carry9 * 19;
+    h9 -= carry9 << 25;
+    carry0 = (h0 + (int64_t) (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;
+}
+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 __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;
+    q = (h0 + q) >> 26;
+    q = (h1 + q) >> 25;
+    q = (h2 + q) >> 26;
+    q = (h3 + q) >> 25;
+    q = (h4 + q) >> 26;
+    q = (h5 + q) >> 25;
+    q = (h6 + q) >> 26;
+    q = (h7 + q) >> 25;
+    q = (h8 + q) >> 26;
+    q = (h9 + q) >> 25;
+    h0 += 19 * q;
+    carry0 = h0 >> 26;
+    h1 += carry0;
+    h0 -= carry0 << 26;
+    carry1 = h1 >> 25;
+    h2 += carry1;
+    h1 -= carry1 << 25;
+    carry2 = h2 >> 26;
+    h3 += carry2;
+    h2 -= carry2 << 26;
+    carry3 = h3 >> 25;
+    h4 += carry3;
+    h3 -= carry3 << 25;
+    carry4 = h4 >> 26;
+    h5 += carry4;
+    h4 -= carry4 << 26;
+    carry5 = h5 >> 25;
+    h6 += carry5;
+    h5 -= carry5 << 25;
+    carry6 = h6 >> 26;
+    h7 += carry6;
+    h6 -= carry6 << 26;
+    carry7 = h7 >> 25;
+    h8 += carry7;
+    h7 -= carry7 << 25;
+    carry8 = h8 >> 26;
+    h9 += carry8;
+    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);
+}

libs/fe.cuh

@@ -0,0 +1,23 @@
+#ifndef __FE_H
+#define __FE_H
+#include <fixedint.h>
+typedef int32_t fe[10];
+void __host__ __device__ fe_0(fe h);
+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);
+#endif

libs/fixedint.h

@@ -0,0 +1,51 @@
+#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 <stdint.h>
+    #define FIXEDINT_H_INCLUDED
+    #if defined(__WATCOMC__) && __WATCOMC__ >= 1250 && !defined(UINT64_C)
+        #include <limits.h>
+        #define UINT64_C(x) (x + (UINT64_MAX - UINT64_MAX))
+    #endif
+#endif
+#ifndef FIXEDINT_H_INCLUDED
+    #define FIXEDINT_H_INCLUDED
+    #include <limits.h>
+    #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

libs/ge.cu

@@ -0,0 +1,343 @@
+#include <ge.cuh>
+#include <precomp_data.h>
+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) {
+    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->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_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) {
+    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) {
+    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) {
+    fe t0;
+    fe_sq(r->X, p->X);
+    fe_sq(r->Z, p->Y);
+    fe_sq2(r->T, p->Z);
+    fe_add(r->Y, p->X, p->Y);
+    fe_sq(t0, r->Y);
+    fe_add(r->Y, r->Z, r->X);
+    fe_sub(r->Z, r->Z, r->X);
+    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) {
+    ge_p2 q;
+    ge_p3_to_p2(&q, p);
+    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;
+    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;
+}
+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;
+    x >>= 63;
+    return (unsigned char) x;
+}
+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) {
+    ge_precomp minust;
+    unsigned char bnegative = negative(b);
+    unsigned char babs = b - (((-bnegative) & b) << 1);
+    fe_1(t->yplusx);
+    fe_1(t->yminusx);
+    fe_0(t->xy2d);
+    cmov(t, &base[pos][0], equal(babs, 1));
+    cmov(t, &base[pos][1], equal(babs, 2));
+    cmov(t, &base[pos][2], equal(babs, 3));
+    cmov(t, &base[pos][3], equal(babs, 4));
+    cmov(t, &base[pos][4], equal(babs, 5));
+    cmov(t, &base[pos][5], equal(babs, 6));
+    cmov(t, &base[pos][6], equal(babs, 7));
+    cmov(t, &base[pos][7], equal(babs, 8));
+    fe_copy(minust.yplusx, t->yminusx);
+    fe_copy(minust.yminusx, t->yplusx);
+    fe_neg(minust.xy2d, t->xy2d);
+    cmov(t, &minust, bnegative);
+}
+void __device__ __host__ ge_scalarmult_base(ge_p3* h, const unsigned char* a) {
+    signed char e[64], carry;
+    ge_p1p1 r;
+    ge_p2 s;
+    ge_precomp t;
+    int i;
+#pragma unroll 32
+    for (i = 0; i < 32; i++) {
+        e[2 * i] = a[i] & 15;
+        e[2 * i + 1] = a[i] >> 4;
+    }
+#pragma unroll
+    for (i = 0, carry = 0; i < 63; i++) {
+        e[i] += carry;
+        carry = (e[i] + 8) >> 4;
+        e[i] -= carry << 4;
+    }
+    e[63] += carry;
+    ge_p3_0(h);
+#pragma unroll 32
+    for (i = 1; i < 64; i += 2) {
+        select(&t, i >> 1, e[i]);
+        ge_madd(&r, h, &t);
+        ge_p1p1_to_p3(h, &r);
+    }
+    ge_p3_dbl(&r, h);
+    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_p2(&s, &r); ge_p2_dbl(&r, &s);
+    ge_p1p1_to_p3(h, &r);
+#pragma unroll 32
+    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;
+}

libs/ge.cuh

@@ -0,0 +1,49 @@
+#ifndef __GE_H
+#define __GE_H
+#include <fe.cuh>
+typedef struct {
+  fe X;
+  fe Y;
+  fe Z;
+} ge_p2;
+typedef struct {
+  fe X;
+  fe Y;
+  fe Z;
+  fe T;
+} ge_p3;
+typedef struct {
+  fe X;
+  fe Y;
+  fe Z;
+  fe T;
+} ge_p1p1;
+typedef struct {
+  fe yplusx;
+  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

sources/main.cu

@@ -105,7 +105,7 @@ __global__ void KeyGenKernel(curandState* randStates) {
         Key32 seed;
         KeysBox32 keys;
         rmbytes(seed, &localState);
-        ed25519_keygen(keys.PrivateKey, keys.PublicKey, seed);
+        ed25519_create_keypair(keys.PrivateKey, keys.PublicKey, seed);
         if (unsigned zeros = getZeros(keys.PublicKey); zeros > atomicMax(&d_high, zeros)) {
             Addr16 raw;
             Key32 inv;
@@ -150,11 +150,11 @@ int main(int argc, char* argv[]) {
         free(h_seeds);
         cudaFree(d_seeds);
         cudaFree(rst);
-        return EXIT_FAILURE;
+        return 1;
     }
     free(h_seeds);
     cudaFree(d_seeds);
-    KeyGenKernel<<<tTh / THREADS_P_B, THREADS_P_B>>>(rst);
+    KeyGenKernel << <tTh / THREADS_P_B, THREADS_P_B >> > (rst);
     cudaFree(rst);
     return 0;
 }