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

fixes

Browse Source
This commit is contained in:
rcxpony
2025-03-18 03:46:01 +05:00
parent 1499fc5d28
commit a9a693dc07
4 changed files with 10 additions and 28 deletions
Download Patch File Download Diff File Expand all files Collapse all files

16
libs/edsign.cu
View File

@@ -10,22 +10,6 @@ __device__ __forceinline__ void expand_key(unsigned char* __restrict__ expanded,
expanded[0] &= 0xf8;
expanded[31] = (expanded[31] & 0x7F) | 0x40;
}
/*
__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];
ed25519_unproject(x, y, &p);
ed25519_pack(r, x, y);
}
__device__ void ed25519_keygen(unsigned char private_key[64], unsigned char public_key[32], unsigned char random_seed[32]) {
unsigned char expanded[64];
expand_key(expanded, random_seed);
sm_pack(public_key, expanded);
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);

4
libs/edsign.cuh
View File

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

2
libs/f25519.cu
View File

@@ -23,7 +23,7 @@ void __device__ __host__ fe_1(fe __restrict__ h) {
h[8] = 0;
h[9] = 0;
}
void __device__ __host__ fe_add(fe h, const fe& __restrict__ f, const fe& __restrict__ g) {
void __device__ __host__ fe_add(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];

16
sources/main.cu
View File

@@ -7,7 +7,8 @@
#include <string.cuh>
#include <keymanip.cuh>
#include <defines.h>
__device__ unsigned d_high = 0x10;
__device__ unsigned d_high = 0x14;
__device__ int parameters(const char* arg) noexcept {
if ((cstring_find(arg, "--altitude") == 0 && cstring_length(arg) == 10) ||
(cstring_find(arg, "-a") == 0 && cstring_length(arg) == 2)) {
@@ -49,14 +50,13 @@ __device__ __forceinline__ unsigned char zeroCounter(unsigned int x) 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) {
unsigned word = (static_cast<unsigned>(v[i]) << 24) | (static_cast<unsigned>(v[i + 1]) << 16) |
(static_cast<unsigned>(v[i + 2]) << 8) | (static_cast<unsigned>(v[i + 3]));
if (word == 0)
#pragma unroll 8
for (unsigned char i = 0; i < 32; i += 4) {
unsigned w = (static_cast<unsigned>(v[i]) << 24) | (static_cast<unsigned>(v[i + 1]) << 16) | (static_cast<unsigned>(v[i + 2]) << 8) | (static_cast<unsigned>(v[i + 3]));
if (w == 0)
leadZeros += 32;
else {
leadZeros += zeroCounter(word);
leadZeros += zeroCounter(w);
break;
}
}
@@ -95,7 +95,7 @@ __device__ __forceinline__ void rmbytes(unsigned char* __restrict__ buf, curandS
buf[i * 4 + 3] = static_cast<unsigned char>((r >> 24) & 0xFF);
}
}
__global__ void KeyGenKernel(curandState* randStates) {
__global__ void KeyGenKernel(curandState* __restrict__ randStates) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
curandState localState = randStates[idx];
#ifdef DEBUG