yggm/libs/ed25519.cuh

#pragma once
#include <f25519.cuh>
#define F25519_SIZE 32
struct ed25519_pt {
    unsigned char x[F25519_SIZE], y[F25519_SIZE], t[F25519_SIZE], z[F25519_SIZE];
};
__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[F25519_SIZE] = {
    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[F25519_SIZE] = {
    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__ __forceinline__ void ed25519_project(struct ed25519_pt* p, const unsigned char* x, const unsigned char* y) {
    f25519_copy(p->x, x);
    f25519_copy(p->y, y);
    f25519_load(p->z, 1);
    f25519_mul__distinct(p->t, x, y);
}
__device__ __forceinline__ void ed25519_unproject(unsigned char* x, unsigned char* y, const struct ed25519_pt* p) {
    unsigned char z1[F25519_SIZE];
    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);
}
__device__ __forceinline__ void ed25519_pack(unsigned char* c, const unsigned char* x, const unsigned char* y) {
    unsigned char tmp[F25519_SIZE];
    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;
}
__device__ __forceinline__ void ed25519_add(struct ed25519_pt* r, const struct ed25519_pt* p1, const struct ed25519_pt* p2) {
    unsigned char a[F25519_SIZE], b[F25519_SIZE], c[F25519_SIZE], d[F25519_SIZE];
    unsigned char e[F25519_SIZE], f[F25519_SIZE], g[F25519_SIZE], h[F25519_SIZE];
    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);
}
__device__ __forceinline__ void ed25519_double(struct ed25519_pt* r, const struct ed25519_pt* p) {
    unsigned char a[F25519_SIZE], b[F25519_SIZE], c[F25519_SIZE];
    unsigned char e[F25519_SIZE], f[F25519_SIZE], g[F25519_SIZE], h[F25519_SIZE];
    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);
}
__device__ __forceinline__ void ed25519_copy(struct ed25519_pt* dst, const struct ed25519_pt* 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);
}
__device__ __forceinline__ void ed25519_smult(struct ed25519_pt* r_out, const struct ed25519_pt* p, const unsigned char* e) {
    struct ed25519_pt r = ed25519_neutral;
    for (int i = 255; i >= 0; i--) {
        struct ed25519_pt s;
        ed25519_double(&r, &r);
        ed25519_add(&s, &r, p);
        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);
    }
    ed25519_copy(r_out, &r);
}
__device__ __forceinline__ void ed25519_prepare(unsigned char* e) {
    e[0] &= 0xf8;
    e[31] &= 0x7f;
    e[31] |= 0x40;
}