rcxpony/yggm
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

optimized

Browse Source
This commit is contained in:
rcxpony
2025-03-17 15:04:47 +05:00
parent bc814d988c
commit 7ea0e62697
14 changed files with 724 additions and 842 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
README.md
View File

@@ -13,7 +13,8 @@ make -j$(nproc)
./yggmcu -t 10
```
# ToDo
- [x] CUDA support (slow)
- [x] CUDA support
- [x] Support for avx2
- [ ] Support for sse4
- [ ] (CUDA) Optimize internal algorithms
#

211
libs/ed25519.cu
View File

@@ -1,98 +1,129 @@
#include <ed25519.cuh>
#include <f25519.cuh>
__device__ __constant__ struct ed25519_pt ed25519_base = {
{0x1a,0xd5,0x25,0x8f,0x60,0x2d,0x56,0xc9,0xb2,0xa7,0x25,0x95,0x60,0xc7,0x2c,0x69,0x5c,0xdc,0xd6,0xfd,0x31,0xe2,0xa4,0xc0,0xfe,0x53,0x6e,0xcd,0xd3,0x36,0x69,0x21},
{0x58,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66,0x66},
{0xa3,0xdd,0xb7,0xa5,0xb3,0x8a,0xde,0x6d,0xf5,0x52,0x51,0x77,0x80,0x9f,0xf0,0x20,0x7d,0xe3,0xab,0x64,0x8e,0x4e,0xea,0x66,0x65,0x76,0x8b,0xd7,0x0f,0x5f,0x87,0x67},
{1,0}
};
__device__ __constant__ struct ed25519_pt ed25519_neutral = {
{0}, {1,0}, {0}, {1,0}
};
__device__ __constant__ unsigned char ed25519_d[32] = {
0xa3,0x78,0x59,0x13,0xca,0x4d,0xeb,0x75,0xab,0xd8,0x41,0x41,0x4d,0x0a,0x70,0x00,
0x98,0xe8,0x79,0x77,0x79,0x40,0xc7,0x8c,0x73,0xfe,0x6f,0x2b,0xee,0x6c,0x03,0x52
};
__device__ __constant__ unsigned char ed25519_k[32] = {
0x59,0xf1,0xb2,0x26,0x94,0x9b,0xd6,0xeb,0x56,0xb1,0x83,0x82,0x9a,0x14,0xe0,0x00,
0x30,0xd1,0xf3,0xee,0xf2,0x80,0x8e,0x19,0xe7,0xfc,0xdf,0x56,0xdc,0xd9,0x06,0x24
};
__device__ void ed25519_unproject(unsigned char* __restrict__ x, unsigned char* __restrict__ y, const struct ed25519_pt* __restrict__ p) {
unsigned char __align__(32) z1[32];
f25519_inv__distinct(z1, p->z);
f25519_mul__distinct(x, p->x, z1);
f25519_mul__distinct(y, p->y, z1);
f25519_normalize(x);
f25519_normalize(y);
#include <precomp_data.h>
void __host__ __device__ ge_madd(ge_p1p1* __restrict__ r, const ge_p3* __restrict__ p, const ge_precomp* __restrict__ 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);
}
__device__ void ed25519_pack(unsigned char* __restrict__ c, const unsigned char* __restrict__ x, const unsigned char* __restrict__ y) {
unsigned char __align__(32) tmp[32];
unsigned char parity;
f25519_copy(tmp, x);
f25519_normalize(tmp);
parity = (tmp[0] & 1) << 7;
f25519_copy(c, y);
f25519_normalize(c);
c[31] |= parity;
void __host__ __device__ ge_p1p1_to_p2(ge_p2* __restrict__ r, const ge_p1p1* __restrict__ 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);
}
__device__ __forceinline__ void ed25519_add(struct ed25519_pt* __restrict__ r, const struct ed25519_pt* __restrict__ p1, const struct ed25519_pt* __restrict__ p2) {
unsigned char __align__(32) a[32], __align__(32) b[32], __align__(32) c[32], __align__(32) d[32], __align__(32) e[32], __align__(32) f[32], __align__(32) g[32], __align__(32) h[32];
f25519_sub(c, p1->y, p1->x);
f25519_sub(d, p2->y, p2->x);
f25519_mul__distinct(a, c, d);
f25519_add(c, p1->y, p1->x);
f25519_add(d, p2->y, p2->x);
f25519_mul__distinct(b, c, d);
f25519_mul__distinct(d, p1->t, p2->t);
f25519_mul__distinct(c, d, ed25519_k);
f25519_mul__distinct(d, p1->z, p2->z);
f25519_add(d, d, d);
f25519_sub(e, b, a);
f25519_sub(f, d, c);
f25519_add(g, d, c);
f25519_add(h, b, a);
f25519_mul__distinct(r->x, e, f);
f25519_mul__distinct(r->y, g, h);
f25519_mul__distinct(r->t, e, h);
f25519_mul__distinct(r->z, f, g);
void __host__ __device__ ge_p1p1_to_p3(ge_p3* __restrict__ r, const ge_p1p1* __restrict__ 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);
}
__device__ __forceinline__ void ed25519_double(struct ed25519_pt* __restrict__ r, const struct ed25519_pt* __restrict__ p) {
unsigned char __align__(32) a[32], __align__(32) b[32], __align__(32) c[32], __align__(32) e[32], __align__(32) f[32], __align__(32) g[32], __align__(32) h[32];
f25519_mul__distinct(a, p->x, p->x);
f25519_mul__distinct(b, p->y, p->y);
f25519_mul__distinct(c, p->z, p->z);
f25519_add(c, c, c);
f25519_add(f, p->x, p->y);
f25519_mul__distinct(e, f, f);
f25519_sub(e, e, a);
f25519_sub(e, e, b);
f25519_sub(g, b, a);
f25519_sub(f, g, c);
f25519_neg(h, b);
f25519_sub(h, h, a);
f25519_mul__distinct(r->x, e, f);
f25519_mul__distinct(r->y, g, h);
f25519_mul__distinct(r->t, e, h);
f25519_mul__distinct(r->z, f, g);
void __host__ __device__ ge_p2_dbl(ge_p1p1* __restrict__ r, const ge_p2* __restrict__ p) {
fe t0;
fe_mul(r->X, p->X, p->X);
fe_mul(r->Z, p->Y, p->Y);
fe_sq2(r->T, p->Z);
fe_add(r->Y, p->X, p->Y);
fe_mul(t0, r->Y, 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);
}
__device__ __forceinline__ void ed25519_copy(struct ed25519_pt* __restrict__ dst, const struct ed25519_pt* __restrict__ src) {
f25519_copy(dst->x, src->x);
f25519_copy(dst->y, src->y);
f25519_copy(dst->t, src->t);
f25519_copy(dst->z, src->z);
void __host__ __device__ ge_p3_dbl(ge_p1p1* __restrict__ r, const ge_p3* __restrict__ p) {
ge_p2 q;
fe_copy(q.X, p->X);
fe_copy(q.Y, p->Y);
fe_copy(q.Z, p->Z);
ge_p2_dbl(r, &q);
}
__device__ void ed25519_smult(ed25519_pt* r_out, const unsigned char* __restrict__ e) {
ed25519_pt r = ed25519_neutral;
#pragma unroll
for (int i = 255; i >= 0; i--) {
struct ed25519_pt s;
ed25519_double(&r, &r);
ed25519_add(&s, &r, &ed25519_base);
unsigned char bit = (e[i >> 3] >> (i & 7)) & 1;
f25519_select(r.x, r.x, s.x, bit);
f25519_select(r.y, r.y, s.y, bit);
f25519_select(r.z, r.z, s.z, bit);
f25519_select(r.t, r.t, s.t, bit);
void ge_p3_tobytes(unsigned char* __restrict__ s, const ge_p3* __restrict__ h) {
fe recip, x, 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 long x = b ^ c;
x -= 1;
x >>= 63;
return (unsigned char)x;
}
static void __host__ __device__ cmov(ge_precomp* __restrict__ t, const ge_precomp* __restrict__ 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 long x = b;
x >>= 63;
unsigned char bnegative = static_cast<unsigned char>(x);
unsigned char babs = b - (((-bnegative) & b) << 1);
fe_1(t->yplusx);
fe_1(t->yminusx);
#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));
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* __restrict__ h, const unsigned char* __restrict__ 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 63
for (i = 0, carry = 0; i < 63; i++) {
e[i] += carry;
carry = (e[i] + 8) >> 4;
e[i] -= carry << 4;
}
e[63] += carry;
#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
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);
}
ed25519_copy(r_out, &r);
}

46
libs/ed25519.cuh
View File

@@ -1,14 +1,34 @@
#ifndef __ED25519_CUH
#define __ED25519_CUH
struct ed25519_pt { unsigned char x[32], y[32], t[32], z[32]; };
extern __device__ __constant__ struct ed25519_pt ed25519_base;
extern __device__ __constant__ struct ed25519_pt ed25519_neutral;
extern __device__ __constant__ unsigned char ed25519_d[32];
extern __device__ __constant__ unsigned char ed25519_k[32];
__device__ void ed25519_unproject(unsigned char* x, unsigned char* y, const struct ed25519_pt* p);
__device__ void ed25519_pack(unsigned char* c, const unsigned char* x, const unsigned char* y);
__device__ __forceinline__ void ed25519_add(struct ed25519_pt* r, const struct ed25519_pt* p1, const struct ed25519_pt* p2);
__device__ __forceinline__ void ed25519_double(struct ed25519_pt* r, const struct ed25519_pt* p);
__device__ __forceinline__ void ed25519_copy(struct ed25519_pt* dst, const struct ed25519_pt* src);
__device__ void ed25519_smult(struct ed25519_pt* r_out, const unsigned char* e);
#ifndef __ED25519_H
#define __ED25519_H
#include <f25519.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;
void __host__ __device__ ge_p3_tobytes(unsigned char *s, const ge_p3 *h);
void __host__ __device__ ge_madd(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_dbl(ge_p1p1 *r, const ge_p2 *p);
void __host__ __device__ ge_p3_dbl(ge_p1p1 *r, const ge_p3 *p);
void __host__ __device__ ge_p3_to_p2(ge_p2 *r, const ge_p3 *p);
#endif

11
libs/edsign.cu
View File

@@ -1,17 +1,17 @@
#include <edsign.cuh>
#include <ed25519.cuh>
#include <sha512.cuh>
#include <ge.cuh>
__device__ __forceinline__ void expand_key(unsigned char* expanded, const unsigned char* secret) {
#include <ed25519.cuh>
__device__ __forceinline__ void expand_key(unsigned char* __restrict__ expanded, const unsigned char* __restrict__ secret) {
struct sha512_state s;
memcpy(&s, &sha512_initial_state, sizeof(s));
sha512_final(&s, secret);
sha512_get(&s, expanded);
expanded[0] &= 0xf8;
expanded[31] &= 0x7f;
expanded[31] |= 0x40;
expanded[31] = (expanded[31] & 0x7F) | 0x40;
}
__device__ __forceinline__ void sm_pack(unsigned char* r, const unsigned char* k) {
/*
__device__ __forceinline__ void sm_pack(unsigned char* __restrict__ r, const unsigned char* __restrict__ k) {
struct ed25519_pt p;
ed25519_smult(&p, k);
unsigned char x[32], y[32];
@@ -25,6 +25,7 @@ __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);

664
libs/f25519.cu
View File

@@ -1,134 +1,550 @@
__device__ void f25519_copy(unsigned char* __restrict__ x, const unsigned char* __restrict__ a) {
const uint4* src = reinterpret_cast<const uint4*>(a);
uint4* dst = reinterpret_cast<uint4*>(x);
dst[0] = src[0];
dst[1] = src[1];
#include <f25519.cuh>
void __host__ __device__ fe_1(fe __restrict__ h) {
h[0] = 1;
h[1] = 0;
h[2] = 0;
h[3] = 0;
h[4] = 0;
h[5] = 0;
h[6] = 0;
h[7] = 0;
h[8] = 0;
h[9] = 0;
}
__device__ void f25519_select(unsigned char* __restrict__ dst, const unsigned char* __restrict__ z, const unsigned char* __restrict__ o, unsigned char cond) {
unsigned int mask = static_cast<unsigned int>(-cond);
const uint4* vZero = reinterpret_cast<const uint4*>(z);
const uint4* vOne = reinterpret_cast<const uint4*>(o);
uint4* vDst = reinterpret_cast<uint4*>(dst);
uint4 res0, res1;
res0.x = (vZero[0].x & ~mask) | (vOne[0].x & mask);
res0.y = (vZero[0].y & ~mask) | (vOne[0].y & mask);
res0.z = (vZero[0].z & ~mask) | (vOne[0].z & mask);
res0.w = (vZero[0].w & ~mask) | (vOne[0].w & mask);
res1.x = (vZero[1].x & ~mask) | (vOne[1].x & mask);
res1.y = (vZero[1].y & ~mask) | (vOne[1].y & mask);
res1.z = (vZero[1].z & ~mask) | (vOne[1].z & mask);
res1.w = (vZero[1].w & ~mask) | (vOne[1].w & mask);
vDst[0] = res0;
vDst[1] = res1;
void __device__ __host__ fe_add(fe h, const fe& __restrict__ f, const fe& __restrict__ g) {
signed int f0 = f[0];
signed int f1 = f[1];
signed int f2 = f[2];
signed int f3 = f[3];
signed int f4 = f[4];
signed int f5 = f[5];
signed int f6 = f[6];
signed int f7 = f[7];
signed int f8 = f[8];
signed int f9 = f[9];
signed int g0 = g[0];
signed int g1 = g[1];
signed int g2 = g[2];
signed int g3 = g[3];
signed int g4 = g[4];
signed int g5 = g[5];
signed int g6 = g[6];
signed int g7 = g[7];
signed int g8 = g[8];
signed int g9 = g[9];
signed int h0 = f0 + g0;
signed int h1 = f1 + g1;
signed int h2 = f2 + g2;
signed int h3 = f3 + g3;
signed int h4 = f4 + g4;
signed int h5 = f5 + g5;
signed int h6 = f6 + g6;
signed int h7 = f7 + g7;
signed int h8 = f8 + g8;
signed int 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;
}
__device__ void f25519_normalize(unsigned char* __restrict__ x) {
__align__(32) unsigned char minusp[32];
unsigned c = (x[31] >> 7) * 19;
x[31] &= 127;
#pragma unroll 32
for (int i = 0; i < 32; i++) {
c += x[i];
x[i] = static_cast<unsigned char>(c);
c >>= 8;
}
c = 19;
#pragma unroll 31
for (int i = 0; i < 31; i++) {
c += x[i];
minusp[i] = static_cast<unsigned char>(c);
c >>= 8;
}
c += x[31] - 128;
minusp[31] = static_cast<unsigned char>(c);
f25519_select(x, minusp, x, static_cast<unsigned char>((c >> 15) & 1));
void __host__ __device__ fe_cmov(fe __restrict__ f, const fe& __restrict__ g, unsigned int b) {
signed int f0 = f[0];
signed int f1 = f[1];
signed int f2 = f[2];
signed int f3 = f[3];
signed int f4 = f[4];
signed int f5 = f[5];
signed int f6 = f[6];
signed int f7 = f[7];
signed int f8 = f[8];
signed int f9 = f[9];
signed int g0 = g[0];
signed int g1 = g[1];
signed int g2 = g[2];
signed int g3 = g[3];
signed int g4 = g[4];
signed int g5 = g[5];
signed int g6 = g[6];
signed int g7 = g[7];
signed int g8 = g[8];
signed int g9 = g[9];
signed int x0 = f0 ^ g0;
signed int x1 = f1 ^ g1;
signed int x2 = f2 ^ g2;
signed int x3 = f3 ^ g3;
signed int x4 = f4 ^ g4;
signed int x5 = f5 ^ g5;
signed int x6 = f6 ^ g6;
signed int x7 = f7 ^ g7;
signed int x8 = f8 ^ g8;
signed int x9 = f9 ^ g9;
b = (unsigned int)(-(int)b);
x0 &= b;
x1 &= b;
x2 &= b;
x3 &= b;
x4 &= b;
x5 &= b;
x6 &= b;
x7 &= b;
x8 &= b;
x9 &= b;
f[0] = f0 ^ x0;
f[1] = f1 ^ x1;
f[2] = f2 ^ x2;
f[3] = f3 ^ x3;
f[4] = f4 ^ x4;
f[5] = f5 ^ x5;
f[6] = f6 ^ x6;
f[7] = f7 ^ x7;
f[8] = f8 ^ x8;
f[9] = f9 ^ x9;
}
__device__ void f25519_add(unsigned char* __restrict__ r, const unsigned char* __restrict__ a, const unsigned char* __restrict__ b) {
unsigned c = 0;
#pragma unroll 32
for (int i = 0; i < 32; i++) {
c = (c >> 8) + static_cast<unsigned>(a[i]) + static_cast<unsigned>(b[i]);
r[i] = static_cast<unsigned char>(c);
}
r[31] &= 127;
c = (c >> 7) * 19;
#pragma unroll 32
for (int i = 0; i < 32; i++) {
c += r[i];
r[i] = static_cast<unsigned char>(c);
c >>= 8;
}
void __host__ __device__ fe_copy(fe __restrict__ h, const fe& __restrict__ f) {
signed int f0 = f[0];
signed int f1 = f[1];
signed int f2 = f[2];
signed int f3 = f[3];
signed int f4 = f[4];
signed int f5 = f[5];
signed int f6 = f[6];
signed int f7 = f[7];
signed int f8 = f[8];
signed int f9 = f[9];
h[0] = f0;
h[1] = f1;
h[2] = f2;
h[3] = f3;
h[4] = f4;
h[5] = f5;
h[6] = f6;
h[7] = f7;
h[8] = f8;
h[9] = f9;
}
__device__ void f25519_sub(unsigned char* __restrict__ r, const unsigned char* __restrict__ a, const unsigned char* __restrict__ b) {
unsigned c = 218;
#pragma unroll 31
for (int i = 0; i < 31; i++) {
c += 65280 + static_cast<unsigned>(a[i]) - static_cast<unsigned>(b[i]);
r[i] = static_cast<unsigned char>(c);
c >>= 8;
}
c += static_cast<unsigned>(a[31]) - static_cast<unsigned>(b[31]);
r[31] = static_cast<unsigned char>(c & 127);
c = (c >> 7) * 19;
#pragma unroll 32
for (int i = 0; i < 32; i++) {
c += r[i];
r[i] = static_cast<unsigned char>(c);
c >>= 8;
}
void __forceinline__ __device__ __host__ fe_sq(fe __restrict__ h, const fe& __restrict__ f) {
fe_mul(h, f, f);
}
__device__ void f25519_neg(unsigned char* __restrict__ r, const unsigned char* __restrict__ a) {
unsigned c = 218;
#pragma unroll 31
for (int i = 0; i < 31; i++) {
c += 65280 - static_cast<unsigned>(a[i]);
r[i] = static_cast<unsigned char>(c);
c >>= 8;
void fe_invert(fe __restrict__ out, const fe& __restrict__ z) {
fe t0, t1, t2, t3;
int i;
fe_sq(t0, z);
fe_sq(t1, t0);
fe_sq(t1, t1);
fe_mul(t1, z, t1);
fe_mul(t0, t0, t1);
fe_sq(t2, t0);
fe_mul(t1, t1, t2);
fe_sq(t2, t1);
#pragma unroll 4
for (i = 1; i < 5; ++i) {
fe_sq(t2, t2);
}
c -= static_cast<unsigned>(a[31]);
r[31] = static_cast<unsigned char>(c & 127);
c = (c >> 7) * 19;
#pragma unroll 32
for (int i = 0; i < 32; i++) {
c += r[i];
r[i] = static_cast<unsigned char>(c);
c >>= 8;
fe_mul(t1, t2, t1);
fe_sq(t2, t1);
#pragma unroll 9
for (i = 1; i < 10; ++i) {
fe_sq(t2, t2);
}
fe_mul(t2, t2, t1);
fe_sq(t3, t2);
#pragma unroll 19
for (i = 1; i < 20; ++i) {
fe_sq(t3, t3);
}
fe_mul(t2, t3, t2);
fe_sq(t2, t2);
#pragma unroll 9
for (i = 1; i < 10; ++i) {
fe_sq(t2, t2);
}
fe_mul(t1, t2, t1);
fe_sq(t2, t1);
#pragma unroll 49
for (i = 1; i < 50; ++i) {
fe_sq(t2, t2);
}
fe_mul(t2, t2, t1);
fe_sq(t3, t2);
#pragma unroll 99
for (i = 1; i < 100; ++i) {
fe_sq(t3, t3);
}
fe_mul(t2, t3, t2);
fe_sq(t2, t2);
#pragma unroll 49
for (i = 1; i < 50; ++i) {
fe_sq(t2, t2);
}
fe_mul(t1, t2, t1);
fe_sq(t1, t1);
#pragma unroll 4
for (i = 1; i < 5; ++i) {
fe_sq(t1, t1);
}
fe_mul(out, t1, t0);
}
__device__ void f25519_mul__distinct(unsigned char* __restrict__ r, const unsigned char* __restrict__ a, const unsigned char* __restrict__ b) {
unsigned c = 0;
#pragma unroll 32
for (int i = 0; i < 32; i++) {
c >>= 8;
for (int j = 0; j <= i; j++)
c += static_cast<unsigned>(a[j]) * static_cast<unsigned>(b[i - j]);
for (int j = i + 1; j < 32; j++)
c += static_cast<unsigned>(a[j]) * static_cast<unsigned>(b[32 + i - j]) * 38;
r[i] = static_cast<unsigned char>(c);
}
r[31] &= 127;
c = (c >> 7) * 19;
#pragma unroll 32
for (int i = 0; i < 32; i++) {
c += r[i];
r[i] = static_cast<unsigned char>(c);
c >>= 8;
}
int __host__ __device__ fe_isnegative(const fe& __restrict__ f) {
unsigned char s[32];
fe_tobytes(s, f);
return s[0] & 1;
}
__device__ void f25519_inv__distinct(unsigned char* __restrict__ r, const unsigned char* __restrict__ x) {
__align__(32) unsigned char s[32];
f25519_mul__distinct(s, x, x);
f25519_mul__distinct(r, s, x);
#pragma unroll 248
for (int i = 0; i < 248; i++) {
f25519_mul__distinct(s, r, r);
f25519_mul__distinct(r, s, x);
}
f25519_mul__distinct(s, r, r);
f25519_mul__distinct(r, s, s);
f25519_mul__distinct(s, r, x);
f25519_mul__distinct(r, s, s);
f25519_mul__distinct(s, r, r);
f25519_mul__distinct(r, s, x);
f25519_mul__distinct(s, r, r);
f25519_mul__distinct(r, s, x);
__device__ __host__ void fe_mul(fe __restrict__ h, const fe& __restrict__ f, const fe& __restrict__ g) {
long 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];
long 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];
long f0g0 = f0 * g0, f0g1 = f0 * g1, f0g2 = f0 * g2, f0g3 = f0 * g3, f0g4 = f0 * g4, f0g5 = f0 * g5, f0g6 = f0 * g6, f0g7 = f0 * g7, f0g8 = f0 * g8, f0g9 = f0 * g9;
long f1g0 = f1 * g0, f1g1_2 = f1 * g1 << 1, f1g2 = f1 * g2, f1g3_2 = f1 * g3 << 1, f1g4 = f1 * g4, f1g5_2 = f1 * g5 << 1, f1g6 = f1 * g6, f1g7_2 = f1 * g7 << 1, f1g8 = f1 * g8, f1g9_38 = f1 * (19 * g9) << 1;
long f2g0 = f2 * g0, f2g1 = f2 * g1, f2g2 = f2 * g2, f2g3 = f2 * g3, f2g4 = f2 * g4, f2g5 = f2 * g5, f2g6 = f2 * g6, f2g7 = f2 * g7, f2g8_19 = f2 * (19 * g8), f2g9_19 = f2 * (19 * g9);
long f3g0 = f3 * g0, f3g1_2 = f3 * g1 << 1, f3g2 = f3 * g2, f3g3_2 = f3 * g3 << 1, f3g4 = f3 * g4, f3g5_2 = f3 * g5 << 1, f3g6 = f3 * g6, f3g7_38 = f3 * (19 * g7) << 1, f3g8_19 = f3 * (19 * g8), f3g9_38 = f3 * (19 * g9) << 1;
long f4g0 = f4 * g0, f4g1 = f4 * g1, f4g2 = f4 * g2, f4g3 = f4 * g3, f4g4 = f4 * g4, f4g5 = f4 * g5, f4g6_19 = f4 * (19 * g6), f4g7_19 = f4 * (19 * g7), f4g8_19 = f4 * (19 * g8), f4g9_19 = f4 * (19 * g9);
long f5g0 = f5 * g0, f5g1_2 = f5 * g1 << 1, f5g2 = f5 * g2, f5g3_2 = f5 * g3 << 1, f5g4 = f5 * g4, f5g5_38 = f5 * (19 * g5) << 1, f5g6_19 = f5 * (19 * g6), f5g7_38 = f5 * (19 * g7) << 1, f5g8_19 = f5 * (19 * g8), f5g9_38 = f5 * (19 * g9) << 1;
long f6g0 = f6 * g0, f6g1 = f6 * g1, f6g2 = f6 * g2, f6g3 = f6 * g3, f6g4_19 = f6 * (19 * g4), f6g5_19 = f6 * (19 * g5), f6g6_19 = f6 * (19 * g6), f6g7_19 = f6 * (19 * g7), f6g8_19 = f6 * (19 * g8), f6g9_19 = f6 * (19 * g9);
long f7g0 = f7 * g0, f7g1_2 = f7 * g1 << 1, f7g2 = f7 * g2, f7g3_38 = f7 * (19 * g3) << 1, f7g4_19 = f7 * (19 * g4), f7g5_38 = f7 * (19 * g5) << 1, f7g6_19 = f7 * (19 * g6), f7g7_38 = f7 * (19 * g7) << 1, f7g8_19 = f7 * (19 * g8), f7g9_38 = f7 * (19 * g9) << 1;
long f8g0 = f8 * g0, f8g1 = f8 * g1, f8g2_19 = f8 * (19 * g2), f8g3_19 = f8 * (19 * g3), f8g4_19 = f8 * (19 * g4), f8g5_19 = f8 * (19 * g5), f8g6_19 = f8 * (19 * g6), f8g7_19 = f8 * (19 * g7), f8g8_19 = f8 * (19 * g8), f8g9_19 = f8 * (19 * g9);
long f9g0 = f9 * g0, f9g1_38 = f9 * (19 * g1) << 1, f9g2_19 = f9 * (19 * g2), f9g3_38 = f9 * (19 * g3) << 1, f9g4_19 = f9 * (19 * g4), f9g5_38 = f9 * (19 * g5) << 1, f9g6_19 = f9 * (19 * g6), f9g7_38 = f9 * (19 * g7) << 1, f9g8_19 = f9 * (19 * g8), f9g9_38 = f9 * (19 * g9) << 1;
long h0 = f0g0 + f1g9_38 + f2g8_19 + f3g7_38 + f4g6_19 + f5g5_38 + f6g4_19 + f7g3_38 + f8g2_19 + f9g1_38;
long h1 = f0g1 + f1g0 + f2g9_19 + f3g8_19 + f4g7_19 + f5g6_19 + f6g5_19 + f7g4_19 + f8g3_19 + f9g2_19;
long h2 = f0g2 + f1g1_2 + f2g0 + f3g9_38 + f4g8_19 + f5g7_38 + f6g6_19 + f7g5_38 + f8g4_19 + f9g3_38;
long h3 = f0g3 + f1g2 + f2g1 + f3g0 + f4g9_19 + f5g8_19 + f6g7_19 + f7g6_19 + f8g5_19 + f9g4_19;
long h4 = f0g4 + f1g3_2 + f2g2 + f3g1_2 + f4g0 + f5g9_38 + f6g8_19 + f7g7_38 + f8g6_19 + f9g5_38;
long h5 = f0g5 + f1g4 + f2g3 + f3g2 + f4g1 + f5g0 + f6g9_19 + f7g8_19 + f8g7_19 + f9g6_19;
long h6 = f0g6 + f1g5_2 + f2g4 + f3g3_2 + f4g2 + f5g1_2 + f6g0 + f7g9_38 + f8g8_19 + f9g7_38;
long h7 = f0g7 + f1g6 + f2g5 + f3g4 + f4g3 + f5g2 + f6g1 + f7g0 + f8g9_19 + f9g8_19;
long h8 = f0g8 + f1g7_2 + f2g6 + f3g5_2 + f4g4 + f5g3_2 + f6g2 + f7g1_2 + f8g0 + f9g9_38;
long h9 = f0g9 + f1g8 + f2g7 + f3g6 + f4g5 + f5g4 + f6g3 + f7g2 + f8g1 + f9g0;
long carry = (h0 + (1L << 25)) >> 26; h1 += carry; h0 -= carry << 26;
carry = (h4 + (1L << 25)) >> 26; h5 += carry; h4 -= carry << 26;
carry = (h1 + (1L << 24)) >> 25; h2 += carry; h1 -= carry << 25;
carry = (h5 + (1L << 24)) >> 25; h6 += carry; h5 -= carry << 25;
carry = (h2 + (1L << 25)) >> 26; h3 += carry; h2 -= carry << 26;
carry = (h6 + (1L << 25)) >> 26; h7 += carry; h6 -= carry << 26;
carry = (h3 + (1L << 24)) >> 25; h4 += carry; h3 -= carry << 25;
carry = (h7 + (1L << 24)) >> 25; h8 += carry; h7 -= carry << 25;
carry = (h4 + (1L << 25)) >> 26; h5 += carry; h4 -= carry << 26;
carry = (h8 + (1L << 25)) >> 26; h9 += carry; h8 -= carry << 26;
carry = (h9 + (1L << 24)) >> 25; h0 += carry * 19; h9 -= carry << 25;
carry = (h0 + (1L << 25)) >> 26; h1 += carry; h0 -= carry << 26;
h[0] = static_cast<int>(h0);
h[1] = static_cast<int>(h1);
h[2] = static_cast<int>(h2);
h[3] = static_cast<int>(h3);
h[4] = static_cast<int>(h4);
h[5] = static_cast<int>(h5);
h[6] = static_cast<int>(h6);
h[7] = static_cast<int>(h7);
h[8] = static_cast<int>(h8);
h[9] = static_cast<int>(h9);
}
void __device__ __host__ fe_neg(fe h, const fe& __restrict__ f) {
signed int f0 = f[0];
signed int f1 = f[1];
signed int f2 = f[2];
signed int f3 = f[3];
signed int f4 = f[4];
signed int f5 = f[5];
signed int f6 = f[6];
signed int f7 = f[7];
signed int f8 = f[8];
signed int f9 = f[9];
signed int h0 = -f0;
signed int h1 = -f1;
signed int h2 = -f2;
signed int h3 = -f3;
signed int h4 = -f4;
signed int h5 = -f5;
signed int h6 = -f6;
signed int h7 = -f7;
signed int h8 = -f8;
signed int 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 __host__ __device__ fe_sq2(fe __restrict__ h, const fe& __restrict__ 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 = f0 << 1;
int f1_2 = f1 << 1;
int f2_2 = f2 << 1;
int f3_2 = f3 << 1;
int f4_2 = f4 << 1;
int f5_2 = f5 << 1;
int f6_2 = f6 << 1;
int f7_2 = f7 << 1;
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 = f0 * static_cast<long>(f0);
long f0f1_2 = f0_2 * static_cast<long>(f1);
long f0f2_2 = f0_2 * static_cast<long>(f2);
long f0f3_2 = f0_2 * static_cast<long>(f3);
long f0f4_2 = f0_2 * static_cast<long>(f4);
long f0f5_2 = f0_2 * static_cast<long>(f5);
long f0f6_2 = f0_2 * static_cast<long>(f6);
long f0f7_2 = f0_2 * static_cast<long>(f7);
long f0f8_2 = f0_2 * static_cast<long>(f8);
long f0f9_2 = f0_2 * static_cast<long>(f9);
long f1f1_2 = f1_2 * static_cast<long>(f1);
long f1f2_2 = f1_2 * static_cast<long>(f2);
long f1f3_4 = f1_2 * static_cast<long>(f3_2);
long f1f4_2 = f1_2 * static_cast<long>(f4);
long f1f5_4 = f1_2 * static_cast<long>(f5_2);
long f1f6_2 = f1_2 * static_cast<long>(f6);
long f1f7_4 = f1_2 * static_cast<long>(f7_2);
long f1f8_2 = f1_2 * static_cast<long>(f8);
long f1f9_76 = f1_2 * static_cast<long>(f9_38);
long f2f2 = f2 * static_cast<long>(f2);
long f2f3_2 = f2_2 * static_cast<long>(f3);
long f2f4_2 = f2_2 * static_cast<long>(f4);
long f2f5_2 = f2_2 * static_cast<long>(f5);
long f2f6_2 = f2_2 * static_cast<long>(f6);
long f2f7_2 = f2_2 * static_cast<long>(f7);
long f2f8_38 = f2_2 * static_cast<long>(f8_19);
long f2f9_38 = f2 * static_cast<long>(f9_38);
long f3f3_2 = f3_2 * static_cast<long>(f3);
long f3f4_2 = f3_2 * static_cast<long>(f4);
long f3f5_4 = f3_2 * static_cast<long>(f5_2);
long f3f6_2 = f3_2 * static_cast<long>(f6);
long f3f7_76 = f3_2 * static_cast<long>(f7_38);
long f3f8_38 = f3_2 * static_cast<long>(f8_19);
long f3f9_76 = f3_2 * static_cast<long>(f9_38);
long f4f4 = f4 * static_cast<long>(f4);
long f4f5_2 = f4_2 * static_cast<long>(f5);
long f4f6_38 = f4_2 * static_cast<long>(f6_19);
long f4f7_38 = f4 * static_cast<long>(f7_38);
long f4f8_38 = f4_2 * static_cast<long>(f8_19);
long f4f9_38 = f4 * static_cast<long>(f9_38);
long f5f5_38 = f5 * static_cast<long>(f5_38);
long f5f6_38 = f5_2 * static_cast<long>(f6_19);
long f5f7_76 = f5_2 * static_cast<long>(f7_38);
long f5f8_38 = f5_2 * static_cast<long>(f8_19);
long f5f9_76 = f5_2 * static_cast<long>(f9_38);
long f6f6_19 = f6 * static_cast<long>(f6_19);
long f6f7_38 = f6 * static_cast<long>(f7_38);
long f6f8_38 = f6_2 * static_cast<long>(f8_19);
long f6f9_38 = f6 * static_cast<long>(f9_38);
long f7f7_38 = f7 * static_cast<long>(f7_38);
long f7f8_38 = f7_2 * static_cast<long>(f8_19);
long f7f9_76 = f7_2 * static_cast<long>(f9_38);
long f8f8_19 = f8 * static_cast<long>(f8_19);
long f8f9_38 = f8 * static_cast<long>(f9_38);
long f9f9_38 = f9 * static_cast<long>(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 <<= 1;
h1 <<= 1;
h2 <<= 1;
h3 <<= 1;
h4 <<= 1;
h5 <<= 1;
h6 <<= 1;
h7 <<= 1;
h8 <<= 1;
h9 <<= 1;
carry0 = (h0 + static_cast<long>(1 << 25)) >> 26;
h1 += carry0;
h0 -= carry0 << 26;
carry4 = (h4 + static_cast<long>(1 << 25)) >> 26;
h5 += carry4;
h4 -= carry4 << 26;
carry1 = (h1 + static_cast<long>(1 << 24)) >> 25;
h2 += carry1;
h1 -= carry1 << 25;
carry5 = (h5 + static_cast<long>(1 << 24)) >> 25;
h6 += carry5;
h5 -= carry5 << 25;
carry2 = (h2 + static_cast<long>(1 << 25)) >> 26;
h3 += carry2;
h2 -= carry2 << 26;
carry6 = (h6 + static_cast<long>(1 << 25)) >> 26;
h7 += carry6;
h6 -= carry6 << 26;
carry3 = (h3 + static_cast<long>(1 << 24)) >> 25;
h4 += carry3;
h3 -= carry3 << 25;
carry7 = (h7 + static_cast<long>(1 << 24)) >> 25;
h8 += carry7;
h7 -= carry7 << 25;
carry4 = (h4 + static_cast<long>(1 << 25)) >> 26;
h5 += carry4;
h4 -= carry4 << 26;
carry8 = (h8 + static_cast<long>(1 << 25)) >> 26;
h9 += carry8;
h8 -= carry8 << 26;
carry9 = (h9 + static_cast<long>(1 << 24)) >> 25;
h0 += carry9 * 19;
h9 -= carry9 << 25;
carry0 = (h0 + static_cast<long>(1 << 25)) >> 26;
h1 += carry0;
h0 -= carry0 << 26;
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 __restrict__ h, const fe& __restrict__ f, const fe& __restrict__ g) {
signed int f0 = f[0];
signed int f1 = f[1];
signed int f2 = f[2];
signed int f3 = f[3];
signed int f4 = f[4];
signed int f5 = f[5];
signed int f6 = f[6];
signed int f7 = f[7];
signed int f8 = f[8];
signed int f9 = f[9];
signed int g0 = g[0];
signed int g1 = g[1];
signed int g2 = g[2];
signed int g3 = g[3];
signed int g4 = g[4];
signed int g5 = g[5];
signed int g6 = g[6];
signed int g7 = g[7];
signed int g8 = g[8];
signed int g9 = g[9];
signed int h0 = f0 - g0;
signed int h1 = f1 - g1;
signed int h2 = f2 - g2;
signed int h3 = f3 - g3;
signed int h4 = f4 - g4;
signed int h5 = f5 - g5;
signed int h6 = f6 - g6;
signed int h7 = f7 - g7;
signed int h8 = f8 - g8;
signed int 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* __restrict__ s, const fe& __restrict__ h) {
int h0 = h[0], h1 = h[1], h2 = h[2], h3 = h[3], h4 = h[4], h5 = h[5], h6 = h[6], h7 = h[7], h8 = h[8], h9 = h[9], q;
int carry0, carry1, carry2, carry3, carry4, carry5, carry6, carry7, carry8, carry9;
q = (19 * h9 + (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] = static_cast<unsigned char>(h0 >> 0);
s[1] = static_cast<unsigned char>(h0 >> 8);
s[2] = static_cast<unsigned char>(h0 >> 16);
s[3] = static_cast<unsigned char>((h0 >> 24) | (h1 << 2));
s[4] = static_cast<unsigned char>(h1 >> 6);
s[5] = static_cast<unsigned char>(h1 >> 14);
s[6] = static_cast<unsigned char>((h1 >> 22) | (h2 << 3));
s[7] = static_cast<unsigned char>(h2 >> 5);
s[8] = static_cast<unsigned char>(h2 >> 13);
s[9] = static_cast<unsigned char>((h2 >> 21) | (h3 << 5));
s[10] = static_cast<unsigned char>(h3 >> 3);
s[11] = static_cast<unsigned char>(h3 >> 11);
s[12] = static_cast<unsigned char>((h3 >> 19) | (h4 << 6));
s[13] = static_cast<unsigned char>(h4 >> 2);
s[14] = static_cast<unsigned char>(h4 >> 10);
s[15] = static_cast<unsigned char>(h4 >> 18);
s[16] = static_cast<unsigned char>(h5 >> 0);
s[17] = static_cast<unsigned char>(h5 >> 8);
s[18] = static_cast<unsigned char>(h5 >> 16);
s[19] = static_cast<unsigned char>((h5 >> 24) | (h6 << 1));
s[20] = static_cast<unsigned char>(h6 >> 7);
s[21] = static_cast<unsigned char>(h6 >> 15);
s[22] = static_cast<unsigned char>((h6 >> 23) | (h7 << 3));
s[23] = static_cast<unsigned char>(h7 >> 5);
s[24] = static_cast<unsigned char>(h7 >> 13);
s[25] = static_cast<unsigned char>((h7 >> 21) | (h8 << 4));
s[26] = static_cast<unsigned char>(h8 >> 4);
s[27] = static_cast<unsigned char>(h8 >> 12);
s[28] = static_cast<unsigned char>((h8 >> 20) | (h9 << 6));
s[29] = static_cast<unsigned char>(h9 >> 2);
s[30] = static_cast<unsigned char>(h9 >> 10);
s[31] = static_cast<unsigned char>(h9 >> 18);
}

26
libs/f25519.cuh
View File

@@ -1,11 +1,15 @@
#ifndef __F25519_CUH
#define __F25519_CUH
__device__ void f25519_copy(unsigned char* __restrict__ x, const unsigned char* __restrict__ a);
__device__ void f25519_select(unsigned char* __restrict__ dst, const unsigned char* __restrict__ zero, const unsigned char* __restrict__ one, unsigned char cond);
__device__ void f25519_normalize(unsigned char* __restrict__ x);
__device__ void f25519_add(unsigned char* __restrict__ r, const unsigned char* __restrict__ a, const unsigned char* __restrict__ b);
__device__ void f25519_sub(unsigned char* __restrict__ r, const unsigned char* __restrict__ a, const unsigned char* __restrict__ b);
__device__ void f25519_neg(unsigned char* __restrict__ r, const unsigned char* __restrict__ a);
__device__ void f25519_mul__distinct(unsigned char* __restrict__ r, const unsigned char* __restrict__ a, const unsigned char* __restrict__ b);
__device__ void f25519_inv__distinct(unsigned char* __restrict__ r, const unsigned char* __restrict__ x);
#endif
#ifndef __FE_H
#define __FE_H
using fe = signed int[10];
void __device__ __host__ fe_1(fe h);
void __device__ __host__ fe_tobytes(unsigned char *s, const fe& h);
void __device__ __host__ fe_copy(fe h, const fe& f);
int __device__ __host__ fe_isnegative(const fe& f);
void __device__ __host__ 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_sq2(fe h, const fe& f);
void __device__ __host__ fe_mul(fe h, const fe& __restrict__ f, const fe& __restrict__ g);
void __device__ __host__ fe_sub(fe h, const fe& f, const fe& g);
#endif

409
libs/fe.cu
View File

@@ -1,409 +0,0 @@
#include <fe.cuh>
void __host__ __device__ fe_1(fe __restrict__ h) {
h[0] = 1;
#pragma unroll 9
for (int i = 1; i < 10; i++) h[i] = 0;
}
void __host__ __device__ fe_add(fe __restrict__ h, const fe& __restrict__ f, const fe& __restrict__ 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) {
#pragma unroll 10
for (int i = 0; i < 10; i++) f[i] ^= -static_cast<int>(b) & (f[i] ^ g[i]);
}
void __host__ __device__ fe_copy(fe __restrict__ h, const fe& __restrict__ f) {
#pragma unroll 10
for (int i = 0; i < 10; i++) h[i] = f[i];
}
void fe_invert(fe __restrict__ out, const fe& __restrict__ 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& __restrict__ f) {
unsigned char s[32];
fe_tobytes(s, f);
return s[0] & 1;
}
__device__ __host__ void fe_mul(fe __restrict__ h, const fe& __restrict__ f, const fe& __restrict__ g) {
long 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];
long 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];
long f0g0 = f0 * g0, f0g1 = f0 * g1, f0g2 = f0 * g2, f0g3 = f0 * g3, f0g4 = f0 * g4, f0g5 = f0 * g5, f0g6 = f0 * g6, f0g7 = f0 * g7, f0g8 = f0 * g8, f0g9 = f0 * g9;
long f1g0 = f1 * g0, f1g1_2 = f1 * g1 << 1, f1g2 = f1 * g2, f1g3_2 = f1 * g3 << 1, f1g4 = f1 * g4, f1g5_2 = f1 * g5 << 1, f1g6 = f1 * g6, f1g7_2 = f1 * g7 << 1, f1g8 = f1 * g8, f1g9_38 = f1 * (19 * g9) << 1;
long f2g0 = f2 * g0, f2g1 = f2 * g1, f2g2 = f2 * g2, f2g3 = f2 * g3, f2g4 = f2 * g4, f2g5 = f2 * g5, f2g6 = f2 * g6, f2g7 = f2 * g7, f2g8_19 = f2 * (19 * g8), f2g9_19 = f2 * (19 * g9);
long f3g0 = f3 * g0, f3g1_2 = f3 * g1 << 1, f3g2 = f3 * g2, f3g3_2 = f3 * g3 << 1, f3g4 = f3 * g4, f3g5_2 = f3 * g5 << 1, f3g6 = f3 * g6, f3g7_38 = f3 * (19 * g7) << 1, f3g8_19 = f3 * (19 * g8), f3g9_38 = f3 * (19 * g9) << 1;
long f4g0 = f4 * g0, f4g1 = f4 * g1, f4g2 = f4 * g2, f4g3 = f4 * g3, f4g4 = f4 * g4, f4g5 = f4 * g5, f4g6_19 = f4 * (19 * g6), f4g7_19 = f4 * (19 * g7), f4g8_19 = f4 * (19 * g8), f4g9_19 = f4 * (19 * g9);
long f5g0 = f5 * g0, f5g1_2 = f5 * g1 << 1, f5g2 = f5 * g2, f5g3_2 = f5 * g3 << 1, f5g4 = f5 * g4, f5g5_38 = f5 * (19 * g5) << 1, f5g6_19 = f5 * (19 * g6), f5g7_38 = f5 * (19 * g7) << 1, f5g8_19 = f5 * (19 * g8), f5g9_38 = f5 * (19 * g9) << 1;
long f6g0 = f6 * g0, f6g1 = f6 * g1, f6g2 = f6 * g2, f6g3 = f6 * g3, f6g4_19 = f6 * (19 * g4), f6g5_19 = f6 * (19 * g5), f6g6_19 = f6 * (19 * g6), f6g7_19 = f6 * (19 * g7), f6g8_19 = f6 * (19 * g8), f6g9_19 = f6 * (19 * g9);
long f7g0 = f7 * g0, f7g1_2 = f7 * g1 << 1, f7g2 = f7 * g2, f7g3_38 = f7 * (19 * g3) << 1, f7g4_19 = f7 * (19 * g4), f7g5_38 = f7 * (19 * g5) << 1, f7g6_19 = f7 * (19 * g6), f7g7_38 = f7 * (19 * g7) << 1, f7g8_19 = f7 * (19 * g8), f7g9_38 = f7 * (19 * g9) << 1;
long f8g0 = f8 * g0, f8g1 = f8 * g1, f8g2_19 = f8 * (19 * g2), f8g3_19 = f8 * (19 * g3), f8g4_19 = f8 * (19 * g4), f8g5_19 = f8 * (19 * g5), f8g6_19 = f8 * (19 * g6), f8g7_19 = f8 * (19 * g7), f8g8_19 = f8 * (19 * g8), f8g9_19 = f8 * (19 * g9);
long f9g0 = f9 * g0, f9g1_38 = f9 * (19 * g1) << 1, f9g2_19 = f9 * (19 * g2), f9g3_38 = f9 * (19 * g3) << 1, f9g4_19 = f9 * (19 * g4), f9g5_38 = f9 * (19 * g5) << 1, f9g6_19 = f9 * (19 * g6), f9g7_38 = f9 * (19 * g7) << 1, f9g8_19 = f9 * (19 * g8), f9g9_38 = f9 * (19 * g9) << 1;
long h0 = f0g0 + f1g9_38 + f2g8_19 + f3g7_38 + f4g6_19 + f5g5_38 + f6g4_19 + f7g3_38 + f8g2_19 + f9g1_38;
long h1 = f0g1 + f1g0 + f2g9_19 + f3g8_19 + f4g7_19 + f5g6_19 + f6g5_19 + f7g4_19 + f8g3_19 + f9g2_19;
long h2 = f0g2 + f1g1_2 + f2g0 + f3g9_38 + f4g8_19 + f5g7_38 + f6g6_19 + f7g5_38 + f8g4_19 + f9g3_38;
long h3 = f0g3 + f1g2 + f2g1 + f3g0 + f4g9_19 + f5g8_19 + f6g7_19 + f7g6_19 + f8g5_19 + f9g4_19;
long h4 = f0g4 + f1g3_2 + f2g2 + f3g1_2 + f4g0 + f5g9_38 + f6g8_19 + f7g7_38 + f8g6_19 + f9g5_38;
long h5 = f0g5 + f1g4 + f2g3 + f3g2 + f4g1 + f5g0 + f6g9_19 + f7g8_19 + f8g7_19 + f9g6_19;
long h6 = f0g6 + f1g5_2 + f2g4 + f3g3_2 + f4g2 + f5g1_2 + f6g0 + f7g9_38 + f8g8_19 + f9g7_38;
long h7 = f0g7 + f1g6 + f2g5 + f3g4 + f4g3 + f5g2 + f6g1 + f7g0 + f8g9_19 + f9g8_19;
long h8 = f0g8 + f1g7_2 + f2g6 + f3g5_2 + f4g4 + f5g3_2 + f6g2 + f7g1_2 + f8g0 + f9g9_38;
long h9 = f0g9 + f1g8 + f2g7 + f3g6 + f4g5 + f5g4 + f6g3 + f7g2 + f8g1 + f9g0;
long carry = (h0 + (1L << 25)) >> 26; h1 += carry; h0 -= carry << 26;
carry = (h4 + (1L << 25)) >> 26; h5 += carry; h4 -= carry << 26;
carry = (h1 + (1L << 24)) >> 25; h2 += carry; h1 -= carry << 25;
carry = (h5 + (1L << 24)) >> 25; h6 += carry; h5 -= carry << 25;
carry = (h2 + (1L << 25)) >> 26; h3 += carry; h2 -= carry << 26;
carry = (h6 + (1L << 25)) >> 26; h7 += carry; h6 -= carry << 26;
carry = (h3 + (1L << 24)) >> 25; h4 += carry; h3 -= carry << 25;
carry = (h7 + (1L << 24)) >> 25; h8 += carry; h7 -= carry << 25;
carry = (h4 + (1L << 25)) >> 26; h5 += carry; h4 -= carry << 26;
carry = (h8 + (1L << 25)) >> 26; h9 += carry; h8 -= carry << 26;
carry = (h9 + (1L << 24)) >> 25; h0 += carry * 19; h9 -= carry << 25;
carry = (h0 + (1L << 25)) >> 26; h1 += carry; h0 -= carry << 26;
h[0] = static_cast<int>(h0);
h[1] = static_cast<int>(h1);
h[2] = static_cast<int>(h2);
h[3] = static_cast<int>(h3);
h[4] = static_cast<int>(h4);
h[5] = static_cast<int>(h5);
h[6] = static_cast<int>(h6);
h[7] = static_cast<int>(h7);
h[8] = static_cast<int>(h8);
h[9] = static_cast<int>(h9);
}
void __device__ __host__ fe_neg(fe h, const fe& f) {
#pragma unroll 10
for (unsigned char x = 0; x < 10; x++) {
h[x] = -f[x];
}
}
void __device__ __host__ fe_sq(fe h, const fe& f) {
fe_mul(h, f, f);
}
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 = f0 << 1;
int f1_2 = f1 << 1;
int f2_2 = f2 << 1;
int f3_2 = f3 << 1;
int f4_2 = f4 << 1;
int f5_2 = f5 << 1;
int f6_2 = f6 << 1;
int f7_2 = f7 << 1;
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 = f0 * static_cast<long>(f0);
long f0f1_2 = f0_2 * static_cast<long>(f1);
long f0f2_2 = f0_2 * static_cast<long>(f2);
long f0f3_2 = f0_2 * static_cast<long>(f3);
long f0f4_2 = f0_2 * static_cast<long>(f4);
long f0f5_2 = f0_2 * static_cast<long>(f5);
long f0f6_2 = f0_2 * static_cast<long>(f6);
long f0f7_2 = f0_2 * static_cast<long>(f7);
long f0f8_2 = f0_2 * static_cast<long>(f8);
long f0f9_2 = f0_2 * static_cast<long>(f9);
long f1f1_2 = f1_2 * static_cast<long>(f1);
long f1f2_2 = f1_2 * static_cast<long>(f2);
long f1f3_4 = f1_2 * static_cast<long>(f3_2);
long f1f4_2 = f1_2 * static_cast<long>(f4);
long f1f5_4 = f1_2 * static_cast<long>(f5_2);
long f1f6_2 = f1_2 * static_cast<long>(f6);
long f1f7_4 = f1_2 * static_cast<long>(f7_2);
long f1f8_2 = f1_2 * static_cast<long>(f8);
long f1f9_76 = f1_2 * static_cast<long>(f9_38);
long f2f2 = f2 * static_cast<long>(f2);
long f2f3_2 = f2_2 * static_cast<long>(f3);
long f2f4_2 = f2_2 * static_cast<long>(f4);
long f2f5_2 = f2_2 * static_cast<long>(f5);
long f2f6_2 = f2_2 * static_cast<long>(f6);
long f2f7_2 = f2_2 * static_cast<long>(f7);
long f2f8_38 = f2_2 * static_cast<long>(f8_19);
long f2f9_38 = f2 * static_cast<long>(f9_38);
long f3f3_2 = f3_2 * static_cast<long>(f3);
long f3f4_2 = f3_2 * static_cast<long>(f4);
long f3f5_4 = f3_2 * static_cast<long>(f5_2);
long f3f6_2 = f3_2 * static_cast<long>(f6);
long f3f7_76 = f3_2 * static_cast<long>(f7_38);
long f3f8_38 = f3_2 * static_cast<long>(f8_19);
long f3f9_76 = f3_2 * static_cast<long>(f9_38);
long f4f4 = f4 * static_cast<long>(f4);
long f4f5_2 = f4_2 * static_cast<long>(f5);
long f4f6_38 = f4_2 * static_cast<long>(f6_19);
long f4f7_38 = f4 * static_cast<long>(f7_38);
long f4f8_38 = f4_2 * static_cast<long>(f8_19);
long f4f9_38 = f4 * static_cast<long>(f9_38);
long f5f5_38 = f5 * static_cast<long>(f5_38);
long f5f6_38 = f5_2 * static_cast<long>(f6_19);
long f5f7_76 = f5_2 * static_cast<long>(f7_38);
long f5f8_38 = f5_2 * static_cast<long>(f8_19);
long f5f9_76 = f5_2 * static_cast<long>(f9_38);
long f6f6_19 = f6 * static_cast<long>(f6_19);
long f6f7_38 = f6 * static_cast<long>(f7_38);
long f6f8_38 = f6_2 * static_cast<long>(f8_19);
long f6f9_38 = f6 * static_cast<long>(f9_38);
long f7f7_38 = f7 * static_cast<long>(f7_38);
long f7f8_38 = f7_2 * static_cast<long>(f8_19);
long f7f9_76 = f7_2 * static_cast<long>(f9_38);
long f8f8_19 = f8 * static_cast<long>(f8_19);
long f8f9_38 = f8 * static_cast<long>(f9_38);
long f9f9_38 = f9 * static_cast<long>(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;
h3 += h3;
h4 += h4;
h5 += h5;
h6 += h6;
h7 += h7;
h8 += h8;
h9 += h9;
carry0 = (h0 + static_cast<long>(1 << 25)) >> 26;
h1 += carry0;
h0 -= carry0 << 26;
carry4 = (h4 + static_cast<long>(1 << 25)) >> 26;
h5 += carry4;
h4 -= carry4 << 26;
carry1 = (h1 + static_cast<long>(1 << 24)) >> 25;
h2 += carry1;
h1 -= carry1 << 25;
carry5 = (h5 + static_cast<long>(1 << 24)) >> 25;
h6 += carry5;
h5 -= carry5 << 25;
carry2 = (h2 + static_cast<long>(1 << 25)) >> 26;
h3 += carry2;
h2 -= carry2 << 26;
carry6 = (h6 + static_cast<long>(1 << 25)) >> 26;
h7 += carry6;
h6 -= carry6 << 26;
carry3 = (h3 + static_cast<long>(1 << 24)) >> 25;
h4 += carry3;
h3 -= carry3 << 25;
carry7 = (h7 + static_cast<long>(1 << 24)) >> 25;
h8 += carry7;
h7 -= carry7 << 25;
carry4 = (h4 + static_cast<long>(1 << 25)) >> 26;
h5 += carry4;
h4 -= carry4 << 26;
carry8 = (h8 + static_cast<long>(1 << 25)) >> 26;
h9 += carry8;
h8 -= carry8 << 26;
carry9 = (h9 + static_cast<long>(1 << 24)) >> 25;
h0 += carry9 * 19;
h9 -= carry9 << 25;
carry0 = (h0 + static_cast<long>(1 << 25)) >> 26;
h1 += carry0;
h0 -= carry0 << 26;
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 __restrict__ h, const fe& __restrict__ f, const fe& __restrict__ g) {
#pragma unroll 10
for (unsigned char x = 0; x < 10; x++) {
h[x] = f[x] - g[x];
}
}
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 + (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);
}

16
libs/fe.cuh
View File

@@ -1,16 +0,0 @@
#ifndef __FE_H
#define __FE_H
using fe = int[10];
void __device__ __host__ fe_1(fe h);
void __device__ __host__ fe_tobytes(unsigned char *s, const fe& h);
void __device__ __host__ fe_copy(fe h, const fe& f);
int __device__ __host__ fe_isnegative(const fe& f);
void __device__ __host__ 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 __device__ __host__ fe_sq2(fe h, const fe& f);
void __device__ __host__ fe_mul(fe h, const fe& __restrict__ f, const fe& __restrict__ g);
void __device__ __host__ fe_sub(fe h, const fe& f, const fe& g);
#endif

129
libs/ge.cu
View File

@@ -1,129 +0,0 @@
#include <ge.cuh>
#include <precomp_data.h>
void __host__ __device__ ge_madd(ge_p1p1* __restrict__ r, const ge_p3* __restrict__ p, const ge_precomp* __restrict__ 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_p1p1_to_p2(ge_p2* __restrict__ r, const ge_p1p1* __restrict__ 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* __restrict__ r, const ge_p1p1* __restrict__ 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_dbl(ge_p1p1* __restrict__ r, const ge_p2* __restrict__ 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_dbl(ge_p1p1* __restrict__ r, const ge_p3* __restrict__ p) {
ge_p2 q;
fe_copy(q.X, p->X);
fe_copy(q.Y, p->Y);
fe_copy(q.Z, p->Z);
ge_p2_dbl(r, &q);
}
void ge_p3_tobytes(unsigned char* __restrict__ s, const ge_p3* __restrict__ h) {
fe recip, x, 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 long x = b ^ c;
x -= 1;
x >>= 63;
return (unsigned char)x;
}
static void __host__ __device__ cmov(ge_precomp* __restrict__ t, const ge_precomp* __restrict__ 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 long x = b;
x >>= 63;
unsigned char bnegative = static_cast<unsigned char>(x);
unsigned char babs = b - (((-bnegative) & b) << 1);
fe_1(t->yplusx);
fe_1(t->yminusx);
#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));
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* __restrict__ h, const unsigned char* __restrict__ 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 63
for (i = 0, carry = 0; i < 63; i++) {
e[i] += carry;
carry = (e[i] + 8) >> 4;
e[i] -= carry << 4;
}
e[63] += carry;
#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
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);
}
}

34
libs/ge.cuh
View File

@@ -1,34 +0,0 @@
#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;
void __host__ __device__ ge_p3_tobytes(unsigned char *s, const ge_p3 *h);
void __host__ __device__ ge_madd(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_dbl(ge_p1p1 *r, const ge_p2 *p);
void __host__ __device__ ge_p3_dbl(ge_p1p1 *r, const ge_p3 *p);
void __host__ __device__ ge_p3_to_p2(ge_p2 *r, const ge_p3 *p);
#endif

5
libs/keymanip.cu
View File

@@ -1,7 +1,7 @@
#include <keymanip.cuh>
static __constant__ const char* hexDigits = "0123456789abcdef";
__device__ ds64 ktos(const unsigned char* key) noexcept {
ds64 str;
const char* hexDigits = "0123456789abcdef";
#pragma unroll 32
for (unsigned char i = 0; i < 32; i++) {
str.data[2 * i] = hexDigits[key[i] >> 4];
@@ -12,11 +12,10 @@ __device__ ds64 ktos(const unsigned char* key) noexcept {
}
__device__ ds46 getAddr(const unsigned char rawAddr[16]) noexcept {
ds46 addrStr;
const char* hexDigits = "0123456789abcdef";
unsigned pos = 0;
#pragma unroll 8
for (unsigned char group = 0; group < 8; group++) {
int idx = group * 2;
int idx = group << 1;
addrStr.data[pos++] = hexDigits[rawAddr[idx] >> 4];
addrStr.data[pos++] = hexDigits[rawAddr[idx] & 0x0F];
addrStr.data[pos++] = hexDigits[rawAddr[idx + 1] >> 4];

3
libs/keymanip.cuh
View File

@@ -11,8 +11,7 @@ using Key32 = unsigned char[32];
struct __align__(32) KeysBox32 {
Key32 PublicKey;
Key32 PrivateKey;
};
__device__ ds64 ktos(const unsigned char* key) noexcept;
};__device__ ds64 ktos(const unsigned char* key) noexcept;
__device__ ds46 getAddr(const unsigned char rawAddr[16]) noexcept;
__device__ void getRawAddress(int lErase, Key32& InvertedPublicKey, Addr16& rawAddr) noexcept;
__device__ void invertKey(const unsigned char* key, unsigned char* inverted);

6
libs/precomp_data.h
View File

@@ -1,5 +1,5 @@
#include <ge.cuh>
__device__ __constant__ ge_precomp Bi[8] = { {
#include <ed25519.cuh>
__device__ __constant__ const ge_precomp Bi[8] = { {
{ 25967493, -14356035, 29566456, 3660896, -12694345, 4014787, 27544626, -11754271, -6079156, 2047605 },
{ -12545711, 934262, -2722910, 3049990, -727428, 9406986, 12720692, 5043384, 19500929, -15469378 },
{ -8738181, 4489570, 9688441, -14785194, 10184609, -12363380, 29287919, 11864899, -24514362, -4438546 },
@@ -40,7 +40,7 @@ __device__ __constant__ ge_precomp Bi[8] = { {
{ -3099351, 10324967, -2241613, 7453183, -5446979, -2735503, -13812022, -16236442, -32461234, -12290683 },
},
};
__device__ __constant__ ge_precomp base[32][8] = {
__device__ __constant__ const ge_precomp base[32][8] = {
{
{
{ 25967493, -14356035, 29566456, 3660896, -12694345, 4014787, 27544626, -11754271, -6079156, 2047605 },

3
sources/main.cu
View File

@@ -106,8 +106,7 @@ __device__ __forceinline__ void rmbytes(unsigned char* buf, curandState* state)
__global__ void KeyGenKernel(curandState* randStates) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
curandState localState = randStates[idx];
int x = 1;
while (x < 0xFFFFFFFF) {
while (true) {
Key32 seed;
KeysBox32 keys;
rmbytes(seed, &localState);