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 05:32:00 +05:00
parent b7240b3d8d
commit bc814d988c
5 changed files with 94 additions and 340 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

388
libs/fe.cu
View File

@ -1,23 +1,22 @@
#include <fe.cuh> #include <fe.cuh>
void __host__ __device__ fe_1(fe h) { void __host__ __device__ fe_1(fe __restrict__ h) {
h[0] = 1; h[0] = 1;
#pragma unroll 10 #pragma unroll 9
for (int i = 1; i < 10; i++) h[i] = 0; for (int i = 1; i < 10; i++) h[i] = 0;
} }
void __host__ __device__ fe_add(int h[10], const fe& f, const fe& g) { void __host__ __device__ fe_add(fe __restrict__ h, const fe& __restrict__ f, const fe& __restrict__ g) {
#pragma unroll 10 #pragma unroll 10
for (int i = 0; i < 10; i++) h[i] = f[i] + g[i]; 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) { void __host__ __device__ fe_cmov(fe f, const fe& g, unsigned int b) {
int mask = -((int)b);
#pragma unroll 10 #pragma unroll 10
for (int i = 0; i < 10; i++) f[i] ^= mask & (f[i] ^ g[i]); for (int i = 0; i < 10; i++) f[i] ^= -static_cast<int>(b) & (f[i] ^ g[i]);
} }
void __host__ __device__ fe_copy(fe h, const fe& f) { void __host__ __device__ fe_copy(fe __restrict__ h, const fe& __restrict__ f) {
#pragma unroll 10 #pragma unroll 10
for (int i = 0; i < 10; i++) h[i] = f[i]; for (int i = 0; i < 10; i++) h[i] = f[i];
} }
void fe_invert(fe out, const fe& z) { void fe_invert(fe __restrict__ out, const fe& __restrict__ z) {
fe t0; fe t0;
fe t1; fe t1;
fe t2; fe t2;
@ -90,302 +89,65 @@ void fe_invert(fe out, const fe& z) {
} }
fe_mul(out, t1, t0); fe_mul(out, t1, t0);
} }
int __host__ __device__ fe_isnegative(const fe& f) { int __host__ __device__ fe_isnegative(const fe& __restrict__ f) {
unsigned char s[32]; unsigned char s[32];
fe_tobytes(s, f); fe_tobytes(s, f);
return s[0] & 1; return s[0] & 1;
} }
__device__ __host__ void fe_mul(fe h, const fe& f, const fe& g) { __device__ __host__ void fe_mul(fe __restrict__ h, const fe& __restrict__ f, const fe& __restrict__ g) {
const long f0g0 = static_cast<long>(f[0]) * static_cast<long>(g[0]); 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];
const long f0g1 = static_cast<long>(f[0]) * static_cast<long>(g[1]); 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];
const long f0g2 = static_cast<long>(f[0]) * static_cast<long>(g[2]); 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;
const long f0g3 = static_cast<long>(f[0]) * static_cast<long>(g[3]); 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;
const long f0g4 = static_cast<long>(f[0]) * static_cast<long>(g[4]); 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);
const long f0g5 = static_cast<long>(f[0]) * static_cast<long>(g[5]); 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;
const long f0g6 = static_cast<long>(f[0]) * static_cast<long>(g[6]); 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);
const long f0g7 = static_cast<long>(f[0]) * static_cast<long>(g[7]); 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;
const long f0g8 = static_cast<long>(f[0]) * static_cast<long>(g[8]); 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);
const long f0g9 = static_cast<long>(f[0]) * static_cast<long>(g[9]); 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;
const long f1g0 = static_cast<long>(f[1]) * static_cast<long>(g[0]); 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);
const long f1g1_2 = static_cast<long>(2 * f[1]) * static_cast<long>(g[1]); 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;
const long f1g2 = static_cast<long>(f[1]) * static_cast<long>(g[2]); long h0 = f0g0 + f1g9_38 + f2g8_19 + f3g7_38 + f4g6_19 + f5g5_38 + f6g4_19 + f7g3_38 + f8g2_19 + f9g1_38;
const long f1g3_2 = static_cast<long>(2 * f[1]) * static_cast<long>(g[3]); long h1 = f0g1 + f1g0 + f2g9_19 + f3g8_19 + f4g7_19 + f5g6_19 + f6g5_19 + f7g4_19 + f8g3_19 + f9g2_19;
const long f1g4 = static_cast<long>(f[1]) * static_cast<long>(g[4]); long h2 = f0g2 + f1g1_2 + f2g0 + f3g9_38 + f4g8_19 + f5g7_38 + f6g6_19 + f7g5_38 + f8g4_19 + f9g3_38;
const long f1g5_2 = static_cast<long>(2 * f[1]) * static_cast<long>(g[5]); long h3 = f0g3 + f1g2 + f2g1 + f3g0 + f4g9_19 + f5g8_19 + f6g7_19 + f7g6_19 + f8g5_19 + f9g4_19;
const long f1g6 = static_cast<long>(f[1]) * static_cast<long>(g[6]); long h4 = f0g4 + f1g3_2 + f2g2 + f3g1_2 + f4g0 + f5g9_38 + f6g8_19 + f7g7_38 + f8g6_19 + f9g5_38;
const long f1g7_2 = static_cast<long>(2 * f[1]) * static_cast<long>(g[7]); long h5 = f0g5 + f1g4 + f2g3 + f3g2 + f4g1 + f5g0 + f6g9_19 + f7g8_19 + f8g7_19 + f9g6_19;
const long f1g8 = static_cast<long>(f[1]) * static_cast<long>(g[8]); long h6 = f0g6 + f1g5_2 + f2g4 + f3g3_2 + f4g2 + f5g1_2 + f6g0 + f7g9_38 + f8g8_19 + f9g7_38;
const long f1g9_38 = static_cast<long>(2 * f[1]) * static_cast<long>(19 * g[9]); long h7 = f0g7 + f1g6 + f2g5 + f3g4 + f4g3 + f5g2 + f6g1 + f7g0 + f8g9_19 + f9g8_19;
const long f2g0 = static_cast<long>(f[2]) * static_cast<long>(g[0]); long h8 = f0g8 + f1g7_2 + f2g6 + f3g5_2 + f4g4 + f5g3_2 + f6g2 + f7g1_2 + f8g0 + f9g9_38;
const long f2g1 = static_cast<long>(f[2]) * static_cast<long>(g[1]); long h9 = f0g9 + f1g8 + f2g7 + f3g6 + f4g5 + f5g4 + f6g3 + f7g2 + f8g1 + f9g0;
const long f2g2 = static_cast<long>(f[2]) * static_cast<long>(g[2]); long carry = (h0 + (1L << 25)) >> 26; h1 += carry; h0 -= carry << 26;
const long f2g3 = static_cast<long>(f[2]) * static_cast<long>(g[3]); carry = (h4 + (1L << 25)) >> 26; h5 += carry; h4 -= carry << 26;
const long f2g4 = static_cast<long>(f[2]) * static_cast<long>(g[4]); carry = (h1 + (1L << 24)) >> 25; h2 += carry; h1 -= carry << 25;
const long f2g5 = static_cast<long>(f[2]) * static_cast<long>(g[5]); carry = (h5 + (1L << 24)) >> 25; h6 += carry; h5 -= carry << 25;
const long f2g6 = static_cast<long>(f[2]) * static_cast<long>(g[6]); carry = (h2 + (1L << 25)) >> 26; h3 += carry; h2 -= carry << 26;
const long f2g7 = static_cast<long>(f[2]) * static_cast<long>(g[7]); carry = (h6 + (1L << 25)) >> 26; h7 += carry; h6 -= carry << 26;
const long f2g8_19 = static_cast<long>(f[2]) * static_cast<long>(19 * g[8]); carry = (h3 + (1L << 24)) >> 25; h4 += carry; h3 -= carry << 25;
const long f2g9_19 = static_cast<long>(f[2]) * static_cast<long>(19 * g[9]); carry = (h7 + (1L << 24)) >> 25; h8 += carry; h7 -= carry << 25;
const long f3g0 = static_cast<long>(f[3]) * static_cast<long>(g[0]); carry = (h4 + (1L << 25)) >> 26; h5 += carry; h4 -= carry << 26;
const long f3g1_2 = static_cast<long>(2 * f[3]) * static_cast<long>(g[1]); carry = (h8 + (1L << 25)) >> 26; h9 += carry; h8 -= carry << 26;
const long f3g2 = static_cast<long>(f[3]) * static_cast<long>(g[2]); carry = (h9 + (1L << 24)) >> 25; h0 += carry * 19; h9 -= carry << 25;
const long f3g3_2 = static_cast<long>(2 * f[3]) * static_cast<long>(g[3]); carry = (h0 + (1L << 25)) >> 26; h1 += carry; h0 -= carry << 26;
const long f3g4 = static_cast<long>(f[3]) * static_cast<long>(g[4]); h[0] = static_cast<int>(h0);
const long f3g5_2 = static_cast<long>(2 * f[3]) * static_cast<long>(g[5]); h[1] = static_cast<int>(h1);
const long f3g6 = static_cast<long>(f[3]) * static_cast<long>(g[6]); h[2] = static_cast<int>(h2);
const long f3g7_38 = static_cast<long>(2 * f[3]) * static_cast<long>(19 * g[7]); h[3] = static_cast<int>(h3);
const long f3g8_19 = static_cast<long>(f[3]) * static_cast<long>(19 * g[8]); h[4] = static_cast<int>(h4);
const long f3g9_38 = static_cast<long>(2 * f[3]) * static_cast<long>(19 * g[9]); h[5] = static_cast<int>(h5);
const long f4g0 = static_cast<long>(f[4]) * static_cast<long>(g[0]); h[6] = static_cast<int>(h6);
const long f4g1 = static_cast<long>(f[4]) * static_cast<long>(g[1]); h[7] = static_cast<int>(h7);
const long f4g2 = static_cast<long>(f[4]) * static_cast<long>(g[2]); h[8] = static_cast<int>(h8);
const long f4g3 = static_cast<long>(f[4]) * static_cast<long>(g[3]); h[9] = static_cast<int>(h9);
const long f4g4 = static_cast<long>(f[4]) * static_cast<long>(g[4]);
const long f4g5 = static_cast<long>(f[4]) * static_cast<long>(g[5]);
const long f4g6_19 = static_cast<long>(f[4]) * static_cast<long>(19 * g[6]);
const long f4g7_19 = static_cast<long>(f[4]) * static_cast<long>(19 * g[7]);
const long f4g8_19 = static_cast<long>(f[4]) * static_cast<long>(19 * g[8]);
const long f4g9_19 = static_cast<long>(f[4]) * static_cast<long>(19 * g[9]);
const long f5g0 = static_cast<long>(f[5]) * static_cast<long>(g[0]);
const long f5g1_2 = static_cast<long>(2 * f[5]) * static_cast<long>(g[1]);
const long f5g2 = static_cast<long>(f[5]) * static_cast<long>(g[2]);
const long f5g3_2 = static_cast<long>(2 * f[5]) * static_cast<long>(g[3]);
const long f5g4 = static_cast<long>(f[5]) * static_cast<long>(g[4]);
const long f5g5_38 = static_cast<long>(2 * f[5]) * static_cast<long>(19 * g[5]);
const long f5g6_19 = static_cast<long>(f[5]) * static_cast<long>(19 * g[6]);
const long f5g7_38 = static_cast<long>(2 * f[5]) * static_cast<long>(19 * g[7]);
const long f5g8_19 = static_cast<long>(f[5]) * static_cast<long>(19 * g[8]);
const long f5g9_38 = static_cast<long>(2 * f[5]) * static_cast<long>(19 * g[9]);
const long f6g0 = static_cast<long>(f[6]) * static_cast<long>(g[0]);
const long f6g1 = static_cast<long>(f[6]) * static_cast<long>(g[1]);
const long f6g2 = static_cast<long>(f[6]) * static_cast<long>(g[2]);
const long f6g3 = static_cast<long>(f[6]) * static_cast<long>(g[3]);
const long f6g4_19 = static_cast<long>(f[6]) * static_cast<long>(19 * g[4]);
const long f6g5_19 = static_cast<long>(f[6]) * static_cast<long>(19 * g[5]);
const long f6g6_19 = static_cast<long>(f[6]) * static_cast<long>(19 * g[6]);
const long f6g7_19 = static_cast<long>(f[6]) * static_cast<long>(19 * g[7]);
const long f6g8_19 = static_cast<long>(f[6]) * static_cast<long>(19 * g[8]);
const long f6g9_19 = static_cast<long>(f[6]) * static_cast<long>(19 * g[9]);
const long f7g0 = static_cast<long>(f[7]) * static_cast<long>(g[0]);
const long f7g1_2 = static_cast<long>(2 * f[7]) * static_cast<long>(g[1]);
const long f7g2 = static_cast<long>(f[7]) * static_cast<long>(g[2]);
const long f7g3_38 = static_cast<long>(2 * f[7]) * static_cast<long>(19 * g[3]);
const long f7g4_19 = static_cast<long>(f[7]) * static_cast<long>(19 * g[4]);
const long f7g5_38 = static_cast<long>(2 * f[7]) * static_cast<long>(19 * g[5]);
const long f7g6_19 = static_cast<long>(f[7]) * static_cast<long>(19 * g[6]);
const long f7g7_38 = static_cast<long>(2 * f[7]) * static_cast<long>(19 * g[7]);
const long f7g8_19 = static_cast<long>(f[7]) * static_cast<long>(19 * g[8]);
const long f7g9_38 = static_cast<long>(2 * f[7]) * static_cast<long>(19 * g[9]);
const long f8g0 = static_cast<long>(f[8]) * static_cast<long>(g[0]);
const long f8g1 = static_cast<long>(f[8]) * static_cast<long>(g[1]);
const long f8g2_19 = static_cast<long>(f[8]) * static_cast<long>(19 * g[2]);
const long f8g3_19 = static_cast<long>(f[8]) * static_cast<long>(19 * g[3]);
const long f8g4_19 = static_cast<long>(f[8]) * static_cast<long>(19 * g[4]);
const long f8g5_19 = static_cast<long>(f[8]) * static_cast<long>(19 * g[5]);
const long f8g6_19 = static_cast<long>(f[8]) * static_cast<long>(19 * g[6]);
const long f8g7_19 = static_cast<long>(f[8]) * static_cast<long>(19 * g[7]);
const long f8g8_19 = static_cast<long>(f[8]) * static_cast<long>(19 * g[8]);
const long f8g9_19 = static_cast<long>(f[8]) * static_cast<long>(19 * g[9]);
const long f9g0 = static_cast<long>(f[9]) * static_cast<long>(g[0]);
const long f9g1_38 = static_cast<long>(2 * f[9]) * static_cast<long>(19 * g[1]);
const long f9g2_19 = static_cast<long>(f[9]) * static_cast<long>(19 * g[2]);
const long f9g3_38 = static_cast<long>(2 * f[9]) * static_cast<long>(19 * g[3]);
const long f9g4_19 = static_cast<long>(f[9]) * static_cast<long>(19 * g[4]);
const long f9g5_38 = static_cast<long>(2 * f[9]) * static_cast<long>(19 * g[5]);
const long f9g6_19 = static_cast<long>(f[9]) * static_cast<long>(19 * g[6]);
const long f9g7_38 = static_cast<long>(2 * f[9]) * static_cast<long>(19 * g[7]);
const long f9g8_19 = static_cast<long>(f[9]) * static_cast<long>(19 * g[8]);
const long f9g9_38 = static_cast<long>(2 * f[9]) * static_cast<long>(19 * g[9]);
long h0_val = f0g0 + f1g9_38 + f2g8_19 + f3g7_38 + f4g6_19 + f5g5_38 + f6g4_19 + f7g3_38 + f8g2_19 + f9g1_38;
long h1_val = f0g1 + f1g0 + f2g9_19 + f3g8_19 + f4g7_19 + f5g6_19 + f6g5_19 + f7g4_19 + f8g3_19 + f9g2_19;
long h2_val = f0g2 + f1g1_2 + f2g0 + f3g9_38 + f4g8_19 + f5g7_38 + f6g6_19 + f7g5_38 + f8g4_19 + f9g3_38;
long h3_val = f0g3 + f1g2 + f2g1 + f3g0 + f4g9_19 + f5g8_19 + f6g7_19 + f7g6_19 + f8g5_19 + f9g4_19;
long h4_val = f0g4 + f1g3_2 + f2g2 + f3g1_2 + f4g0 + f5g9_38 + f6g8_19 + f7g7_38 + f8g6_19 + f9g5_38;
long h5_val = f0g5 + f1g4 + f2g3 + f3g2 + f4g1 + f5g0 + f6g9_19 + f7g8_19 + f8g7_19 + f9g6_19;
long h6_val = f0g6 + f1g5_2 + f2g4 + f3g3_2 + f4g2 + f5g1_2 + f6g0 + f7g9_38 + f8g8_19 + f9g7_38;
long h7_val = f0g7 + f1g6 + f2g5 + f3g4 + f4g3 + f5g2 + f6g1 + f7g0 + f8g9_19 + f9g8_19;
long h8_val = f0g8 + f1g7_2 + f2g6 + f3g5_2 + f4g4 + f5g3_2 + f6g2 + f7g1_2 + f8g0 + f9g9_38;
long h9_val = f0g9 + f1g8 + f2g7 + f3g6 + f4g5 + f5g4 + f6g3 + f7g2 + f8g1 + f9g0;
long carry;
carry = (h0_val + (static_cast<long>(1LL) << 25)) >> 26;
h1_val += carry; h0_val -= carry << 26;
carry = (h4_val + (static_cast<long>(1LL) << 25)) >> 26;
h5_val += carry; h4_val -= carry << 26;
carry = (h1_val + (static_cast<long>(1LL) << 24)) >> 25;
h2_val += carry; h1_val -= carry << 25;
carry = (h5_val + (static_cast<long>(1LL) << 24)) >> 25;
h6_val += carry; h5_val -= carry << 25;
carry = (h2_val + (static_cast<long>(1LL) << 25)) >> 26;
h3_val += carry; h2_val -= carry << 26;
carry = (h6_val + (static_cast<long>(1LL) << 25)) >> 26;
h7_val += carry; h6_val -= carry << 26;
carry = (h3_val + (static_cast<long>(1LL) << 24)) >> 25;
h4_val += carry; h3_val -= carry << 25;
carry = (h7_val + (static_cast<long>(1LL) << 24)) >> 25;
h8_val += carry; h7_val -= carry << 25;
carry = (h4_val + (static_cast<long>(1LL) << 25)) >> 26;
h5_val += carry; h4_val -= carry << 26;
carry = (h8_val + (static_cast<long>(1LL) << 25)) >> 26;
h9_val += carry; h8_val -= carry << 26;
carry = (h9_val + (static_cast<long>(1LL) << 24)) >> 25;
h0_val += carry * 19; h9_val -= carry << 25;
carry = (h0_val + (static_cast<long>(1LL) << 25)) >> 26;
h1_val += carry; h0_val -= carry << 26;
h[0] = static_cast<int>(h0_val);
h[1] = static_cast<int>(h1_val);
h[2] = static_cast<int>(h2_val);
h[3] = static_cast<int>(h3_val);
h[4] = static_cast<int>(h4_val);
h[5] = static_cast<int>(h5_val);
h[6] = static_cast<int>(h6_val);
h[7] = static_cast<int>(h7_val);
h[8] = static_cast<int>(h8_val);
h[9] = static_cast<int>(h9_val);
} }
void __device__ __host__ fe_neg(fe h, const fe& f) { void __device__ __host__ fe_neg(fe h, const fe& f) {
h[0] = -f[0]; #pragma unroll 10
h[1] = -f[1]; for (unsigned char x = 0; x < 10; x++) {
h[2] = -f[2]; h[x] = -f[x];
h[3] = -f[3]; }
h[4] = -f[4];
h[5] = -f[5];
h[6] = -f[6];
h[7] = -f[7];
h[8] = -f[8];
h[9] = -f[9];
} }
void __device__ __host__ fe_sq(fe h, const fe& f) { void __device__ __host__ fe_sq(fe h, const fe& f) {
int f2 = f[2]; fe_mul(h, f, f);
int f3 = f[3];
int f4 = f[4];
int f5 = f[5];
int f6 = f[6];
int f7 = f[7];
int f8 = f[8];
int f9 = f[9];
int f2_2 = 2 * f2;
int f3_2 = 2 * f3;
int f4_2 = 2 * f4;
int f5_2 = 2 * f5;
int f6_2 = 2 * f6;
int f7_2 = 2 * f7;
int f5_38 = 38 * f5;
int f6_19 = 19 * f6;
int f7_38 = 38 * f7;
int f8_19 = 19 * f8;
int f9_38 = 38 * f9;
long f0f0 = f[0] * static_cast<long>(f[0]);
long f0f1_2 = 2 * f[0] *static_cast<long>(f[1]);
long f0f2_2 = 2 * f[0] * static_cast<long>(f[2]);
long f0f3_2 = 2 * f[0] * static_cast<long>(f[3]);
long f0f4_2 = 2 * f[0] * static_cast<long>(f[4]);
long f0f5_2 = 2 * f[0] * static_cast<long>(f[5]);
long f0f6_2 = 2 * f[0] * static_cast<long>(f[6]);
long f0f7_2 = 2 * f[0] * static_cast<long>(f[7]);
long f0f8_2 = 2 * f[0] * static_cast<long>(f[8]);
long f0f9_2 = 2 * f[0] * static_cast<long>(f[9]);
long f1f1_2 = 2 * f[1] * static_cast<long>(f[1]);
long f1f2_2 = 2 * f[1] * static_cast<long>(f[2]);
long f1f3_4 = 2 * f[1] * static_cast<long>(f[3] * 2);
long f1f4_2 = 2 * f[1] * static_cast<long>(f[4]);
long f1f5_4 = 2 * f[1] * static_cast<long>(f[5] * 2);
long f1f6_2 = 2 * f[1] * static_cast<long>(f[6]);
long f1f7_4 = 2 * f[1] * static_cast<long>(f[7] * 2);
long f1f8_2 = 2 * f[1] * static_cast<long>(f[8]);
long f1f9_76 = 2 * f[1] * static_cast<long>(f[9] * 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, carry1, carry2, carry3, carry4, carry5, carry6, carry7, carry8, carry9;
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 __host__ __device__ fe_sq2(fe h, const fe& f) { void __host__ __device__ fe_sq2(fe h, const fe& f) {
int f0 = f[0]; int f0 = f[0];
@ -398,20 +160,20 @@ void __host__ __device__ fe_sq2(fe h, const fe& f) {
int f7 = f[7]; int f7 = f[7];
int f8 = f[8]; int f8 = f[8];
int f9 = f[9]; int f9 = f[9];
int f0_2 = 2 * f0; int f0_2 = f0 << 1;
int f1_2 = 2 * f1; int f1_2 = f1 << 1;
int f2_2 = 2 * f2; int f2_2 = f2 << 1;
int f3_2 = 2 * f3; int f3_2 = f3 << 1;
int f4_2 = 2 * f4; int f4_2 = f4 << 1;
int f5_2 = 2 * f5; int f5_2 = f5 << 1;
int f6_2 = 2 * f6; int f6_2 = f6 << 1;
int f7_2 = 2 * f7; int f7_2 = f7 << 1;
int f5_38 = 38 * f5; int f5_38 = 38 * f5;
int f6_19 = 19 * f6; int f6_19 = 19 * f6;
int f7_38 = 38 * f7; int f7_38 = 38 * f7;
int f8_19 = 19 * f8; int f8_19 = 19 * f8;
int f9_38 = 38 * f9; int f9_38 = 38 * f9;
long f0f0 = f[0] * static_cast<long>(f[0]); long f0f0 = f0 * static_cast<long>(f0);
long f0f1_2 = f0_2 * static_cast<long>(f1); long f0f1_2 = f0_2 * static_cast<long>(f1);
long f0f2_2 = f0_2 * static_cast<long>(f2); long f0f2_2 = f0_2 * static_cast<long>(f2);
long f0f3_2 = f0_2 * static_cast<long>(f3); long f0f3_2 = f0_2 * static_cast<long>(f3);
@ -543,17 +305,11 @@ void __host__ __device__ fe_sq2(fe h, const fe& f) {
h[8] = (int)h8; h[8] = (int)h8;
h[9] = (int)h9; h[9] = (int)h9;
} }
void __device__ __host__ fe_sub(fe h, const fe& f, const fe& g) { void __device__ __host__ fe_sub(fe __restrict__ h, const fe& __restrict__ f, const fe& __restrict__ g) {
h[0] = (f[0] - g[0]); #pragma unroll 10
h[1] = (f[1] - g[1]); for (unsigned char x = 0; x < 10; x++) {
h[2] = (f[2] - g[2]); h[x] = f[x] - g[x];
h[3] = (f[3] - g[3]); }
h[4] = (f[4] - g[4]);
h[5] = (f[5] - g[5]);
h[6] = (f[6] - g[6]);
h[7] = (f[7] - g[7]);
h[8] = (f[8] - g[8]);
h[9] = (f[9] - g[9]);
} }
void __host__ __device__ fe_tobytes(unsigned char* s, const fe& h) { void __host__ __device__ fe_tobytes(unsigned char* s, const fe& h) {
int h0 = h[0]; int h0 = h[0];
@ -577,7 +333,7 @@ void __host__ __device__ fe_tobytes(unsigned char* s, const fe& h) {
int carry7; int carry7;
int carry8; int carry8;
int carry9; int carry9;
q = (19 * h9 + (((int)1) << 24)) >> 25; q = (19 * h9 + (1 << 24)) >> 25;
q = (h0 + q) >> 26; q = (h0 + q) >> 26;
q = (h1 + q) >> 25; q = (h1 + q) >> 25;
q = (h2 + q) >> 26; q = (h2 + q) >> 26;

10
libs/fe.cuh
View File

@ -3,14 +3,14 @@
using fe = int[10]; using fe = int[10];
void __device__ __host__ fe_1(fe h); void __device__ __host__ fe_1(fe h);
void __device__ __host__ fe_tobytes(unsigned char *s, const fe& h); void __device__ __host__ fe_tobytes(unsigned char *s, const fe& h);
void __host__ __device__ fe_copy(fe h, const fe& f); void __device__ __host__ fe_copy(fe h, const fe& f);
int __host__ __device__ fe_isnegative(const fe& f); int __device__ __host__ fe_isnegative(const fe& f);
void __host__ __device__ fe_cmov(fe f, const fe& g, unsigned int b); 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_neg(fe h, const fe& f);
void __device__ __host__ fe_add(fe h, const fe& f, const fe& g); 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_invert(fe out, const fe& z);
void __device__ __host__ fe_sq(fe h, const fe& f); void __device__ __host__ fe_sq(fe h, const fe& f);
void __host__ __device__ fe_sq2(fe h, const fe& f); void __device__ __host__ fe_sq2(fe h, const fe& f);
void __device__ __host__ fe_mul(fe h, const fe& f, const fe& g); 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); void __device__ __host__ fe_sub(fe h, const fe& f, const fe& g);
#endif #endif

22
libs/ge.cu
View File

@ -1,6 +1,6 @@
#include <ge.cuh> #include <ge.cuh>
#include <precomp_data.h> #include <precomp_data.h>
void __host__ __device__ ge_madd(ge_p1p1* r, const ge_p3* p, const ge_precomp* q) { void __host__ __device__ ge_madd(ge_p1p1* __restrict__ r, const ge_p3* __restrict__ p, const ge_precomp* __restrict__ q) {
fe t0; fe t0;
fe_add(r->X, p->Y, p->X); fe_add(r->X, p->Y, p->X);
fe_sub(r->Y, p->Y, p->X); fe_sub(r->Y, p->Y, p->X);
@ -13,18 +13,18 @@ void __host__ __device__ ge_madd(ge_p1p1* r, const ge_p3* p, const ge_precomp* q
fe_add(r->Z, t0, r->T); fe_add(r->Z, t0, r->T);
fe_sub(r->T, t0, r->T); fe_sub(r->T, t0, r->T);
} }
void __host__ __device__ ge_p1p1_to_p2(ge_p2* r, const ge_p1p1* p) { void __host__ __device__ ge_p1p1_to_p2(ge_p2* __restrict__ r, const ge_p1p1* __restrict__ p) {
fe_mul(r->X, p->X, p->T); fe_mul(r->X, p->X, p->T);
fe_mul(r->Y, p->Y, p->Z); fe_mul(r->Y, p->Y, p->Z);
fe_mul(r->Z, p->Z, p->T); fe_mul(r->Z, p->Z, p->T);
} }
void __host__ __device__ ge_p1p1_to_p3(ge_p3* r, const ge_p1p1* p) { void __host__ __device__ ge_p1p1_to_p3(ge_p3* __restrict__ r, const ge_p1p1* __restrict__ p) {
fe_mul(r->X, p->X, p->T); fe_mul(r->X, p->X, p->T);
fe_mul(r->Y, p->Y, p->Z); fe_mul(r->Y, p->Y, p->Z);
fe_mul(r->Z, p->Z, p->T); fe_mul(r->Z, p->Z, p->T);
fe_mul(r->T, p->X, p->Y); fe_mul(r->T, p->X, p->Y);
} }
void __host__ __device__ ge_p2_dbl(ge_p1p1* r, const ge_p2* p) { void __host__ __device__ ge_p2_dbl(ge_p1p1* __restrict__ r, const ge_p2* __restrict__ p) {
fe t0; fe t0;
fe_sq(r->X, p->X); fe_sq(r->X, p->X);
fe_sq(r->Z, p->Y); fe_sq(r->Z, p->Y);
@ -36,14 +36,14 @@ void __host__ __device__ ge_p2_dbl(ge_p1p1* r, const ge_p2* p) {
fe_sub(r->X, t0, r->Y); fe_sub(r->X, t0, r->Y);
fe_sub(r->T, r->T, r->Z); fe_sub(r->T, r->T, r->Z);
} }
void __host__ __device__ ge_p3_dbl(ge_p1p1* r, const ge_p3* p) { void __host__ __device__ ge_p3_dbl(ge_p1p1* __restrict__ r, const ge_p3* __restrict__ p) {
ge_p2 q; ge_p2 q;
fe_copy(q.X, p->X); fe_copy(q.X, p->X);
fe_copy(q.Y, p->Y); fe_copy(q.Y, p->Y);
fe_copy(q.Z, p->Z); fe_copy(q.Z, p->Z);
ge_p2_dbl(r, &q); ge_p2_dbl(r, &q);
} }
void ge_p3_tobytes(unsigned char* s, const ge_p3* h) { void ge_p3_tobytes(unsigned char* __restrict__ s, const ge_p3* __restrict__ h) {
fe recip, x, y; fe recip, x, y;
fe_invert(recip, h->Z); fe_invert(recip, h->Z);
fe_mul(x, h->X, recip); fe_mul(x, h->X, recip);
@ -57,7 +57,7 @@ static unsigned char __host__ __device__ equal(signed char b, signed char c) {
x >>= 63; x >>= 63;
return (unsigned char)x; return (unsigned char)x;
} }
static void __host__ __device__ cmov(ge_precomp* t, const ge_precomp* u, unsigned char b) { static void __host__ __device__ cmov(ge_precomp* __restrict__ t, const ge_precomp* __restrict__ u, unsigned char b) {
fe_cmov(t->yplusx, u->yplusx, b); fe_cmov(t->yplusx, u->yplusx, b);
fe_cmov(t->yminusx, u->yminusx, b); fe_cmov(t->yminusx, u->yminusx, b);
fe_cmov(t->xy2d, u->xy2d, b); fe_cmov(t->xy2d, u->xy2d, b);
@ -85,7 +85,7 @@ static void __host__ __device__ select(ge_precomp* t, int pos, signed char b) {
fe_neg(minust.xy2d, t->xy2d); fe_neg(minust.xy2d, t->xy2d);
cmov(t, &minust, bnegative); cmov(t, &minust, bnegative);
} }
void __device__ __host__ ge_scalarmult_base(ge_p3* h, const unsigned char* a) { void __device__ __host__ ge_scalarmult_base(ge_p3* __restrict__ h, const unsigned char* __restrict__ a) {
signed char e[64], carry; signed char e[64], carry;
ge_p1p1 r; ge_p1p1 r;
ge_p2 s; ge_p2 s;
@ -96,7 +96,7 @@ void __device__ __host__ ge_scalarmult_base(ge_p3* h, const unsigned char* a) {
e[2 * i] = a[i] & 15; e[2 * i] = a[i] & 15;
e[2 * i + 1] = a[i] >> 4; e[2 * i + 1] = a[i] >> 4;
} }
#pragma unroll #pragma unroll 63
for (i = 0, carry = 0; i < 63; i++) { for (i = 0, carry = 0; i < 63; i++) {
e[i] += carry; e[i] += carry;
carry = (e[i] + 8) >> 4; carry = (e[i] + 8) >> 4;
@ -109,7 +109,7 @@ void __device__ __host__ ge_scalarmult_base(ge_p3* h, const unsigned char* a) {
fe_1(h->Z); fe_1(h->Z);
#pragma unroll 10 #pragma unroll 10
for (int i = 0; i < 10; i++) h->T[i] = 0; for (int i = 0; i < 10; i++) h->T[i] = 0;
#pragma unroll 64 #pragma unroll
for (i = 1; i < 64; i += 2) { for (i = 1; i < 64; i += 2) {
select(&t, i >> 1, e[i]); select(&t, i >> 1, e[i]);
ge_madd(&r, h, &t); ge_madd(&r, h, &t);
@ -120,7 +120,7 @@ void __device__ __host__ ge_scalarmult_base(ge_p3* h, const unsigned char* a) {
ge_p1p1_to_p2(&s, &r); ge_p2_dbl(&r, &s); ge_p1p1_to_p2(&s, &r); ge_p2_dbl(&r, &s);
ge_p1p1_to_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); ge_p1p1_to_p3(h, &r);
#pragma unroll 64 #pragma unroll 32
for (i = 0; i < 64; i += 2) { for (i = 0; i < 64; i += 2) {
select(&t, i >> 1, e[i]); select(&t, i >> 1, e[i]);
ge_madd(&r, h, &t); ge_madd(&r, h, &t);

1
sources/main.cpp
View File

@ -122,7 +122,6 @@ inline void sign_keypair(unsigned char* __restrict pk, unsigned char* __restrict
_mm256_storeu_si256(reinterpret_cast<__m256i*>(sk), _mm256_loadu_si256(reinterpret_cast<const __m256i*>(seed))); _mm256_storeu_si256(reinterpret_cast<__m256i*>(sk), _mm256_loadu_si256(reinterpret_cast<const __m256i*>(seed)));
_mm256_storeu_si256(reinterpret_cast<__m256i*>(sk + 32), _mm256_loadu_si256(reinterpret_cast<const __m256i*>(pk))); _mm256_storeu_si256(reinterpret_cast<__m256i*>(sk + 32), _mm256_loadu_si256(reinterpret_cast<const __m256i*>(pk)));
} }
void miner_thread() noexcept { void miner_thread() noexcept {
alignas(32) Key inv; alignas(32) Key inv;
alignas(32) Key seed; alignas(32) Key seed;

13
sources/main.cu
View File

@ -6,7 +6,7 @@
#include <edsign.cuh> #include <edsign.cuh>
#include <string.cuh> #include <string.cuh>
#include <keymanip.cuh> #include <keymanip.cuh>
#ifndef DEBUG #ifdef RELEASE
#define THREADSPB 256 #define THREADSPB 256
#define THDIVTHPB (tTh / THREADSPB) #define THDIVTHPB (tTh / THREADSPB)
#else #else
@ -68,7 +68,6 @@ __device__ __forceinline__ unsigned char getZeros(const unsigned char* v) noexce
} }
return leadZeros; return leadZeros;
} }
/* /*
__global__ void initRandSeed(curandState* states, const unsigned long seed) { __global__ void initRandSeed(curandState* states, const unsigned long seed) {
int idx = blockIdx.x * blockDim.x + threadIdx.x; int idx = blockIdx.x * blockDim.x + threadIdx.x;
@ -78,14 +77,13 @@ __global__ void initRandSeed(curandState* states, const unsigned long seed) {
__global__ void initRand(curandState* rs, unsigned int* d_seeds) { __global__ void initRand(curandState* rs, unsigned int* d_seeds) {
int id = blockIdx.x * blockDim.x + threadIdx.x; int id = blockIdx.x * blockDim.x + threadIdx.x;
curand_init(clock64() + id * 7919ULL, id, 0, &rs[id]); curand_init(clock64() + id * 7919ULL, id, 0, &rs[id]);
#pragma unroll 10
for (int i = 0; i < 10; i++) { for (int i = 0; i < 10; i++) {
curand(&rs[id]); curand(&rs[id]);
} }
unsigned int seed = curand(&rs[id]); unsigned seed = curand(&rs[id]);
d_seeds[id] = seed; d_seeds[id] = seed;
} }
int checkSeeds(unsigned int* seeds, int count) { int checkSeeds(unsigned* seeds, int count) {
for (int i = 0; i < count; i++) { for (int i = 0; i < count; i++) {
for (int j = i + 1; j < count; j++) { for (int j = i + 1; j < count; j++) {
if (seeds[i] == seeds[j]) { if (seeds[i] == seeds[j]) {
@ -108,7 +106,8 @@ __device__ __forceinline__ void rmbytes(unsigned char* buf, curandState* state)
__global__ void KeyGenKernel(curandState* randStates) { __global__ void KeyGenKernel(curandState* randStates) {
int idx = blockIdx.x * blockDim.x + threadIdx.x; int idx = blockIdx.x * blockDim.x + threadIdx.x;
curandState localState = randStates[idx]; curandState localState = randStates[idx];
while (true) { int x = 1;
while (x < 0xFFFFFFFF) {
Key32 seed; Key32 seed;
KeysBox32 keys; KeysBox32 keys;
rmbytes(seed, &localState); rmbytes(seed, &localState);
@ -122,7 +121,7 @@ __global__ void KeyGenKernel(curandState* randStates) {
d_high = zeros; d_high = zeros;
} }
#ifdef DEBUG #ifdef DEBUG
if ((x & 0xFF) == 0) { if ((++x & 0xFF) == 0) {
printf("Iters: %d\n", x); printf("Iters: %d\n", x);
} }
#endif #endif