fixes

libs/defines.h

@@ -0,0 +1,12 @@
+#ifndef __DEFINES_H
+#define __DEFINES_H
+#ifdef DEBUG
+#define THREADSPB 1
+#define THDIVTHPB 1
+#define WHCOND x < 0xFFFFFFFF
+#else
+#define THREADSPB 256
+#define THDIVTHPB (tTh / THREADSPB)
+#define WHCOND true
+#endif
+#endif

libs/ed25519.cu

@@ -68,8 +68,13 @@ static void __host__ __device__ select(ge_precomp* __restrict__ t, int pos, sign
     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);
+    t->yplusx[0] = 1;
+    t->yminusx[0] = 1;
+#pragma unroll 9
+    for (unsigned char i = 1; i < 10; i++) {
+        t->yplusx[i] = 0;
+        t->yminusx[i] = 0;
+    }
 #pragma unroll 10
     for (int i = 0; i < 10; i++) t->xy2d[i] = 0;
     cmov(t, &base[pos][0], equal(babs, 1));

libs/ed25519.cuh

@@ -1,5 +1,5 @@
-#ifndef __ED25519_H
-#define __ED25519_H
+#ifndef __ED25519_CUH
+#define __ED25519_CUH
 typedef struct {
   signed int X[10], Y[10], Z[10];
 } ge_p2;

libs/f25519.cu

@@ -1,16 +1,4 @@
 #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;
-}
 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];

libs/f25519.cuh

@@ -1,5 +1,5 @@
-#ifndef __FE_H
-#define __FE_H
+#ifndef __F25519_CUH
+#define __F25519_CUH
 using fe = signed int[10];
 void __device__ __host__ fe_1(fe h);
 void __device__ __host__ fe_tobytes(unsigned char *s, const fe& h);

sources/main.cu

@@ -6,13 +6,7 @@
 #include <edsign.cuh>
 #include <string.cuh>
 #include <keymanip.cuh>
-#ifdef RELEASE
-    #define THREADSPB 256
-    #define THDIVTHPB (tTh / THREADSPB)
-#else
-    #define THREADSPB 1
-    #define THDIVTHPB 1
-#endif
+#include <defines.h>
 __device__ unsigned d_high = 0x10;
 __device__ int parameters(const char* arg) noexcept {
     if ((cstring_find(arg, "--altitude") == 0 && cstring_length(arg) == 10) ||
@@ -53,7 +47,7 @@ __global__ void args(char** argv, int argc, int* result) {
 __device__ __forceinline__ unsigned char zeroCounter(unsigned int x) noexcept {
     return x ? static_cast<unsigned char>(__clz(x)) : 32;
 }
-__device__ __forceinline__ unsigned char getZeros(const unsigned char* v) noexcept {
+__device__ __forceinline__ unsigned char getZeros(const unsigned char* __restrict__ v) noexcept {
     unsigned char leadZeros = 0;
 #pragma unroll
     for (int i = 0; i < 32; i += 4) {
@@ -68,7 +62,7 @@ __device__ __forceinline__ unsigned char getZeros(const unsigned char* v) noexce
     }
     return leadZeros;
 }
-__global__ void initRandSeed(curandState* states, const unsigned long seed) {
+__global__ void initRandSeed(curandState* __restrict__ states, const unsigned long seed) {
     int idx = blockIdx.x * blockDim.x + threadIdx.x;
     curand_init(seed, idx, 0, &states[idx]);
 }
@@ -81,7 +75,7 @@ __global__ void initRand(curandState* rs, unsigned int* d_seeds) {
     unsigned seed = curand(&rs[id]);
     d_seeds[id] = seed;
 }
-int checkSeeds(unsigned* seeds, int count) {
+int checkSeeds(unsigned* __restrict__ seeds, int count) {
     for (int i = 0; i < count; i++) {
         for (int j = i + 1; j < count; j++) {
             if (seeds[i] == seeds[j]) {
@@ -91,7 +85,7 @@ int checkSeeds(unsigned* seeds, int count) {
     }
     return 0;
 }
-__device__ __forceinline__ void rmbytes(unsigned char* buf, curandState* state) {
+__device__ __forceinline__ void rmbytes(unsigned char* __restrict__ buf, curandState* __restrict__ state) {
 #pragma unroll 8
     for (int i = 0; i < 8; i++) {
         unsigned r = curand(state);
@@ -104,8 +98,11 @@ __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 = 0;
-    while (x < 0xFFFFFFFF) {
+#ifdef DEBUG 
+    unsigned x = 0;
+#endif
+    
+    while (WHCOND) {
         Key32 seed;
         KeysBox32 keys;
         rmbytes(seed, &localState);