From b54f606d6497309705b31bbc205c6ac1af006885 Mon Sep 17 00:00:00 2001 From: blm Date: Sun, 15 Mar 2020 00:00:11 +0100 Subject: [PATCH 01/10] Shaping readme.md --- README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/README.md b/README.md index 73bebe17..1190d4a2 100644 --- a/README.md +++ b/README.md @@ -127,7 +127,7 @@ There are 3 parameters that affect performance: blocks, threads per block, and k `threads:` The number of threads in a block. This must be a multiple of 32. The default is 256. `Keys per thread:` The number of keys each thread will process. The performance (keys per second) -increases asymptotically with this value. The default is256. Increasing this value will cause the +increases asymptotically with this value. The default is 256. Increasing this value will cause the kernel to run longer, but more keys will be processed. From 9f4ae74541283164a4e2a19f5622075483c90a1c Mon Sep 17 00:00:00 2001 From: blm Date: Wed, 18 Mar 2020 23:23:02 +0100 Subject: [PATCH 02/10] Change global memory access type to vectorized (from scalar) --- CudaKeySearchDevice/CudaDeviceKeys.cu | 32 +++++++--- cudaMath/secp256k1.cuh | 86 +++++++++++++++++---------- 2 files changed, 76 insertions(+), 42 deletions(-) diff --git a/CudaKeySearchDevice/CudaDeviceKeys.cu b/CudaKeySearchDevice/CudaDeviceKeys.cu index d98dbaa7..85db71b5 100644 --- a/CudaKeySearchDevice/CudaDeviceKeys.cu +++ b/CudaKeySearchDevice/CudaDeviceKeys.cu @@ -22,7 +22,7 @@ __device__ unsigned int *ec::getYPtr() return _yPtr[0]; } -__global__ void multiplyStepKernel(const unsigned int *privateKeys, int pointsPerThread, int step, unsigned int *chain, const unsigned int *gxPtr, const unsigned int *gyPtr); +__global__ void multiplyStepKernel(unsigned int *privateKeys, int pointsPerThread, int step, unsigned int *chain, const unsigned int *gxPtr, const unsigned int *gyPtr); int CudaDeviceKeys::getIndex(int block, int thread, int idx) @@ -45,15 +45,22 @@ void CudaDeviceKeys::splatBigInt(unsigned int *dest, int block, int thread, int i.exportWords(value, 8, secp256k1::uint256::BigEndian); int totalThreads = _blocks * _threads; - int threadId = block * _threads + thread; + int threadId = block * _threads * 4 + thread * 4; int base = idx * _blocks * _threads * 8; int index = base + threadId; - for(int k = 0; k < 8; k++) { + for(int k = 0; k < 4; k++) { + dest[index] = value[k]; + index++; + } + + index = base + totalThreads * 4 + threadId; + + for(int k = 4; k < 8; k++) { dest[index] = value[k]; - index += totalThreads; + index++; } } @@ -62,15 +69,22 @@ secp256k1::uint256 CudaDeviceKeys::readBigInt(unsigned int *src, int block, int unsigned int value[8] = { 0 }; int totalThreads = _blocks * _threads; - int threadId = block * _threads + thread; + int threadId = block * _threads * 4 + thread * 4; int base = idx * _blocks * _threads * 8; int index = base + threadId; - for(int k = 0; k < 8; k++) { + for(int k = 0; k < 4; k++) { value[k] = src[index]; - index += totalThreads; + index++; + } + + index = base + totalThreads * 4 + threadId; + + for(int k = 4; k < 8; k++) { + value[k] = src[index]; + index++; } secp256k1::uint256 v(value, secp256k1::uint256::BigEndian); @@ -286,7 +300,7 @@ cudaError_t CudaDeviceKeys::doStep() return err; } -__global__ void multiplyStepKernel(const unsigned int *privateKeys, int pointsPerThread, int step, unsigned int *chain, const unsigned int *gxPtr, const unsigned int *gyPtr) +__global__ void multiplyStepKernel(unsigned int *privateKeys, int pointsPerThread, int step, unsigned int *chain, const unsigned int *gxPtr, const unsigned int *gyPtr) { unsigned int *xPtr = ec::getXPtr(); @@ -394,4 +408,4 @@ bool CudaDeviceKeys::selfTest(const std::vector &privateKeys } return true; -} \ No newline at end of file +} diff --git a/cudaMath/secp256k1.cuh b/cudaMath/secp256k1.cuh index 88a3fed5..75eba4ed 100644 --- a/cudaMath/secp256k1.cuh +++ b/cudaMath/secp256k1.cuh @@ -81,33 +81,44 @@ __device__ static bool equal(const unsigned int *a, const unsigned int *b) /** * Reads an 8-word big integer from device memory */ -__device__ static void readInt(const unsigned int *ara, int idx, unsigned int x[8]) +__device__ static void readInt(unsigned int *ara, int idx, unsigned int x[8]) { - int totalThreads = gridDim.x * blockDim.x; - - int base = idx * totalThreads * 8; - - int threadId = blockDim.x * blockIdx.x + threadIdx.x; - - int index = base + threadId; - - for (int i = 0; i < 8; i++) { - x[i] = ara[index]; - index += totalThreads; - } + uint4 *araTmp = reinterpret_cast(ara); + + unsigned int totalThreads = gridDim.x * blockDim.x; + unsigned int base = idx * totalThreads * 2; + unsigned int threadId = blockDim.x * blockIdx.x + threadIdx.x; + unsigned int index = base + threadId; + + uint4 xTmp = araTmp[index]; + x[0] = xTmp.x; + x[1] = xTmp.y; + x[2] = xTmp.z; + x[3] = xTmp.w; + + index += totalThreads; + + xTmp = araTmp[index]; + x[4] = xTmp.x; + x[5] = xTmp.y; + x[6] = xTmp.z; + x[7] = xTmp.w; } -__device__ static unsigned int readIntLSW(const unsigned int *ara, int idx) +__device__ static unsigned int readIntLSW(unsigned int *ara, int idx) { - int totalThreads = gridDim.x * blockDim.x; + uint4 *araTmp = reinterpret_cast(ara); - int base = idx * totalThreads * 8; + unsigned int totalThreads = gridDim.x * blockDim.x; + unsigned int base = idx * totalThreads * 2; + unsigned int threadId = blockDim.x * blockIdx.x + threadIdx.x; + unsigned int index = base + threadId; - int threadId = blockDim.x * blockIdx.x + threadIdx.x; + index += totalThreads; - int index = base + threadId; + uint4 xTmp = araTmp[index]; - return ara[index + totalThreads * 7]; + return xTmp.w; } /** @@ -115,18 +126,27 @@ __device__ static unsigned int readIntLSW(const unsigned int *ara, int idx) */ __device__ static void writeInt(unsigned int *ara, int idx, const unsigned int x[8]) { - int totalThreads = gridDim.x * blockDim.x; - - int base = idx * totalThreads * 8; - - int threadId = blockDim.x * blockIdx.x + threadIdx.x; - - int index = base + threadId; - - for (int i = 0; i < 8; i++) { - ara[index] = x[i]; - index += totalThreads; - } + uint4 *araTmp = reinterpret_cast(ara); + + unsigned int totalThreads = gridDim.x * blockDim.x; + unsigned int base = idx * totalThreads * 2; + unsigned int threadId = blockDim.x * blockIdx.x + threadIdx.x; + unsigned int index = base + threadId; + + uint4 xTmp; + xTmp.x = x[0]; + xTmp.y = x[1]; + xTmp.z = x[2]; + xTmp.w = x[3]; + araTmp[index] = xTmp; + + index += totalThreads; + + xTmp.x = x[4]; + xTmp.y = x[5]; + xTmp.z = x[6]; + xTmp.w = x[7]; + araTmp[index] = xTmp; } /** @@ -672,7 +692,7 @@ __device__ __forceinline__ static void beginBatchAddWithDouble(const unsigned in writeInt(chain, batchIdx, inverse); } -__device__ static void completeBatchAddWithDouble(const unsigned int *px, const unsigned int *py, const unsigned int *xPtr, const unsigned int *yPtr, int i, int batchIdx, unsigned int *chain, unsigned int *inverse, unsigned int newX[8], unsigned int newY[8]) +__device__ static void completeBatchAddWithDouble(const unsigned int *px, const unsigned int *py, unsigned int *xPtr, unsigned int *yPtr, int i, int batchIdx, unsigned int *chain, unsigned int *inverse, unsigned int newX[8], unsigned int newY[8]) { unsigned int s[8]; unsigned int x[8]; @@ -799,4 +819,4 @@ __device__ __forceinline__ static void doBatchInverse(unsigned int inverse[8]) invModP(inverse); } -#endif \ No newline at end of file +#endif From b5f0fc99e702bf6abbc0055f6c7a4128ab1c790e Mon Sep 17 00:00:00 2001 From: blm Date: Wed, 18 Mar 2020 23:47:07 +0100 Subject: [PATCH 03/10] Update for SM61 --- Makefile | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Makefile b/Makefile index 63758304..9e696e1d 100644 --- a/Makefile +++ b/Makefile @@ -13,7 +13,7 @@ CXX=g++ CXXFLAGS=-O2 -std=c++11 # CUDA variables -COMPUTE_CAP=30 +COMPUTE_CAP=61 NVCC=nvcc NVCCFLAGS=-std=c++11 -gencode=arch=compute_${COMPUTE_CAP},code=\"sm_${COMPUTE_CAP}\" -Xptxas="-v" -Xcompiler "${CXXFLAGS}" CUDA_HOME=/usr/local/cuda From d44220fbfe86fddafacc3760a5433fd4408b5921 Mon Sep 17 00:00:00 2001 From: blm Date: Wed, 21 Apr 2021 22:51:06 +0200 Subject: [PATCH 04/10] Update round fonction name to prevent "c" error with cuda library --- cudaMath/sha256.cuh | 388 ++++++++++++++++++++++---------------------- 1 file changed, 194 insertions(+), 194 deletions(-) diff --git a/cudaMath/sha256.cuh b/cudaMath/sha256.cuh index b04aa9a0..f3a956d7 100644 --- a/cudaMath/sha256.cuh +++ b/cudaMath/sha256.cuh @@ -56,7 +56,7 @@ __device__ __forceinline__ unsigned int s1(unsigned int x) } -__device__ __forceinline__ void round(unsigned int a, unsigned int b, unsigned int c, unsigned int &d, unsigned e, unsigned int f, unsigned int g, unsigned int &h, unsigned int m, unsigned int k) +__device__ __forceinline__ void roundSha256(unsigned int a, unsigned int b, unsigned int c, unsigned int &d, unsigned e, unsigned int f, unsigned int g, unsigned int &h, unsigned int m, unsigned int k) { unsigned int s = CH(e, f, g) + (rotr(e, 6) ^ rotr(e, 11) ^ rotr(e, 25)) + k + m; @@ -98,22 +98,22 @@ __device__ void sha256PublicKey(const unsigned int x[8], const unsigned int y[8] h = _IV[7]; - round(a, b, c, d, e, f, g, h, w[0], _K[0]); - round(h, a, b, c, d, e, f, g, w[1], _K[1]); - round(g, h, a, b, c, d, e, f, w[2], _K[2]); - round(f, g, h, a, b, c, d, e, w[3], _K[3]); - round(e, f, g, h, a, b, c, d, w[4], _K[4]); - round(d, e, f, g, h, a, b, c, w[5], _K[5]); - round(c, d, e, f, g, h, a, b, w[6], _K[6]); - round(b, c, d, e, f, g, h, a, w[7], _K[7]); - round(a, b, c, d, e, f, g, h, w[8], _K[8]); - round(h, a, b, c, d, e, f, g, w[9], _K[9]); - round(g, h, a, b, c, d, e, f, w[10], _K[10]); - round(f, g, h, a, b, c, d, e, w[11], _K[11]); - round(e, f, g, h, a, b, c, d, w[12], _K[12]); - round(d, e, f, g, h, a, b, c, w[13], _K[13]); - round(c, d, e, f, g, h, a, b, w[14], _K[14]); - round(b, c, d, e, f, g, h, a, w[15], _K[15]); + roundSha256(a, b, c, d, e, f, g, h, w[0], _K[0]); + roundSha256(h, a, b, c, d, e, f, g, w[1], _K[1]); + roundSha256(g, h, a, b, c, d, e, f, w[2], _K[2]); + roundSha256(f, g, h, a, b, c, d, e, w[3], _K[3]); + roundSha256(e, f, g, h, a, b, c, d, w[4], _K[4]); + roundSha256(d, e, f, g, h, a, b, c, w[5], _K[5]); + roundSha256(c, d, e, f, g, h, a, b, w[6], _K[6]); + roundSha256(b, c, d, e, f, g, h, a, w[7], _K[7]); + roundSha256(a, b, c, d, e, f, g, h, w[8], _K[8]); + roundSha256(h, a, b, c, d, e, f, g, w[9], _K[9]); + roundSha256(g, h, a, b, c, d, e, f, w[10], _K[10]); + roundSha256(f, g, h, a, b, c, d, e, w[11], _K[11]); + roundSha256(e, f, g, h, a, b, c, d, w[12], _K[12]); + roundSha256(d, e, f, g, h, a, b, c, w[13], _K[13]); + roundSha256(c, d, e, f, g, h, a, b, w[14], _K[14]); + roundSha256(b, c, d, e, f, g, h, a, w[15], _K[15]); @@ -134,22 +134,22 @@ __device__ void sha256PublicKey(const unsigned int x[8], const unsigned int y[8] w[14] = w[14] + s0(w[15]) + w[7] + s1(w[12]); w[15] = w[15] + s0(w[0]) + w[8] + s1(w[13]); - round(a, b, c, d, e, f, g, h, w[0], _K[16]); - round(h, a, b, c, d, e, f, g, w[1], _K[17]); - round(g, h, a, b, c, d, e, f, w[2], _K[18]); - round(f, g, h, a, b, c, d, e, w[3], _K[19]); - round(e, f, g, h, a, b, c, d, w[4], _K[20]); - round(d, e, f, g, h, a, b, c, w[5], _K[21]); - round(c, d, e, f, g, h, a, b, w[6], _K[22]); - round(b, c, d, e, f, g, h, a, w[7], _K[23]); - round(a, b, c, d, e, f, g, h, w[8], _K[24]); - round(h, a, b, c, d, e, f, g, w[9], _K[25]); - round(g, h, a, b, c, d, e, f, w[10], _K[26]); - round(f, g, h, a, b, c, d, e, w[11], _K[27]); - round(e, f, g, h, a, b, c, d, w[12], _K[28]); - round(d, e, f, g, h, a, b, c, w[13], _K[29]); - round(c, d, e, f, g, h, a, b, w[14], _K[30]); - round(b, c, d, e, f, g, h, a, w[15], _K[31]); + roundSha256(a, b, c, d, e, f, g, h, w[0], _K[16]); + roundSha256(h, a, b, c, d, e, f, g, w[1], _K[17]); + roundSha256(g, h, a, b, c, d, e, f, w[2], _K[18]); + roundSha256(f, g, h, a, b, c, d, e, w[3], _K[19]); + roundSha256(e, f, g, h, a, b, c, d, w[4], _K[20]); + roundSha256(d, e, f, g, h, a, b, c, w[5], _K[21]); + roundSha256(c, d, e, f, g, h, a, b, w[6], _K[22]); + roundSha256(b, c, d, e, f, g, h, a, w[7], _K[23]); + roundSha256(a, b, c, d, e, f, g, h, w[8], _K[24]); + roundSha256(h, a, b, c, d, e, f, g, w[9], _K[25]); + roundSha256(g, h, a, b, c, d, e, f, w[10], _K[26]); + roundSha256(f, g, h, a, b, c, d, e, w[11], _K[27]); + roundSha256(e, f, g, h, a, b, c, d, w[12], _K[28]); + roundSha256(d, e, f, g, h, a, b, c, w[13], _K[29]); + roundSha256(c, d, e, f, g, h, a, b, w[14], _K[30]); + roundSha256(b, c, d, e, f, g, h, a, w[15], _K[31]); w[0] = w[0] + s0(w[1]) + w[9] + s1(w[14]); @@ -169,22 +169,22 @@ __device__ void sha256PublicKey(const unsigned int x[8], const unsigned int y[8] w[14] = w[14] + s0(w[15]) + w[7] + s1(w[12]); w[15] = w[15] + s0(w[0]) + w[8] + s1(w[13]); - round(a, b, c, d, e, f, g, h, w[0], _K[32]); - round(h, a, b, c, d, e, f, g, w[1], _K[33]); - round(g, h, a, b, c, d, e, f, w[2], _K[34]); - round(f, g, h, a, b, c, d, e, w[3], _K[35]); - round(e, f, g, h, a, b, c, d, w[4], _K[36]); - round(d, e, f, g, h, a, b, c, w[5], _K[37]); - round(c, d, e, f, g, h, a, b, w[6], _K[38]); - round(b, c, d, e, f, g, h, a, w[7], _K[39]); - round(a, b, c, d, e, f, g, h, w[8], _K[40]); - round(h, a, b, c, d, e, f, g, w[9], _K[41]); - round(g, h, a, b, c, d, e, f, w[10], _K[42]); - round(f, g, h, a, b, c, d, e, w[11], _K[43]); - round(e, f, g, h, a, b, c, d, w[12], _K[44]); - round(d, e, f, g, h, a, b, c, w[13], _K[45]); - round(c, d, e, f, g, h, a, b, w[14], _K[46]); - round(b, c, d, e, f, g, h, a, w[15], _K[47]); + roundSha256(a, b, c, d, e, f, g, h, w[0], _K[32]); + roundSha256(h, a, b, c, d, e, f, g, w[1], _K[33]); + roundSha256(g, h, a, b, c, d, e, f, w[2], _K[34]); + roundSha256(f, g, h, a, b, c, d, e, w[3], _K[35]); + roundSha256(e, f, g, h, a, b, c, d, w[4], _K[36]); + roundSha256(d, e, f, g, h, a, b, c, w[5], _K[37]); + roundSha256(c, d, e, f, g, h, a, b, w[6], _K[38]); + roundSha256(b, c, d, e, f, g, h, a, w[7], _K[39]); + roundSha256(a, b, c, d, e, f, g, h, w[8], _K[40]); + roundSha256(h, a, b, c, d, e, f, g, w[9], _K[41]); + roundSha256(g, h, a, b, c, d, e, f, w[10], _K[42]); + roundSha256(f, g, h, a, b, c, d, e, w[11], _K[43]); + roundSha256(e, f, g, h, a, b, c, d, w[12], _K[44]); + roundSha256(d, e, f, g, h, a, b, c, w[13], _K[45]); + roundSha256(c, d, e, f, g, h, a, b, w[14], _K[46]); + roundSha256(b, c, d, e, f, g, h, a, w[15], _K[47]); @@ -206,22 +206,22 @@ __device__ void sha256PublicKey(const unsigned int x[8], const unsigned int y[8] w[14] = w[14] + s0(w[15]) + w[7] + s1(w[12]); w[15] = w[15] + s0(w[0]) + w[8] + s1(w[13]); - round(a, b, c, d, e, f, g, h, w[0], _K[48]); - round(h, a, b, c, d, e, f, g, w[1], _K[49]); - round(g, h, a, b, c, d, e, f, w[2], _K[50]); - round(f, g, h, a, b, c, d, e, w[3], _K[51]); - round(e, f, g, h, a, b, c, d, w[4], _K[52]); - round(d, e, f, g, h, a, b, c, w[5], _K[53]); - round(c, d, e, f, g, h, a, b, w[6], _K[54]); - round(b, c, d, e, f, g, h, a, w[7], _K[55]); - round(a, b, c, d, e, f, g, h, w[8], _K[56]); - round(h, a, b, c, d, e, f, g, w[9], _K[57]); - round(g, h, a, b, c, d, e, f, w[10], _K[58]); - round(f, g, h, a, b, c, d, e, w[11], _K[59]); - round(e, f, g, h, a, b, c, d, w[12], _K[60]); - round(d, e, f, g, h, a, b, c, w[13], _K[61]); - round(c, d, e, f, g, h, a, b, w[14], _K[62]); - round(b, c, d, e, f, g, h, a, w[15], _K[63]); + roundSha256(a, b, c, d, e, f, g, h, w[0], _K[48]); + roundSha256(h, a, b, c, d, e, f, g, w[1], _K[49]); + roundSha256(g, h, a, b, c, d, e, f, w[2], _K[50]); + roundSha256(f, g, h, a, b, c, d, e, w[3], _K[51]); + roundSha256(e, f, g, h, a, b, c, d, w[4], _K[52]); + roundSha256(d, e, f, g, h, a, b, c, w[5], _K[53]); + roundSha256(c, d, e, f, g, h, a, b, w[6], _K[54]); + roundSha256(b, c, d, e, f, g, h, a, w[7], _K[55]); + roundSha256(a, b, c, d, e, f, g, h, w[8], _K[56]); + roundSha256(h, a, b, c, d, e, f, g, w[9], _K[57]); + roundSha256(g, h, a, b, c, d, e, f, w[10], _K[58]); + roundSha256(f, g, h, a, b, c, d, e, w[11], _K[59]); + roundSha256(e, f, g, h, a, b, c, d, w[12], _K[60]); + roundSha256(d, e, f, g, h, a, b, c, w[13], _K[61]); + roundSha256(c, d, e, f, g, h, a, b, w[14], _K[62]); + roundSha256(b, c, d, e, f, g, h, a, w[15], _K[63]); a += _IV[0]; b += _IV[1]; @@ -246,22 +246,22 @@ __device__ void sha256PublicKey(const unsigned int x[8], const unsigned int y[8] w[0] = (y[7] << 24) | 0x00800000; w[15] = 65 * 8; - round(a, b, c, d, e, f, g, h, w[0], _K[0]); - round(h, a, b, c, d, e, f, g, 0, _K[1]); - round(g, h, a, b, c, d, e, f, 0, _K[2]); - round(f, g, h, a, b, c, d, e, 0, _K[3]); - round(e, f, g, h, a, b, c, d, 0, _K[4]); - round(d, e, f, g, h, a, b, c, 0, _K[5]); - round(c, d, e, f, g, h, a, b, 0, _K[6]); - round(b, c, d, e, f, g, h, a, 0, _K[7]); - round(a, b, c, d, e, f, g, h, 0, _K[8]); - round(h, a, b, c, d, e, f, g, 0, _K[9]); - round(g, h, a, b, c, d, e, f, 0, _K[10]); - round(f, g, h, a, b, c, d, e, 0, _K[11]); - round(e, f, g, h, a, b, c, d, 0, _K[12]); - round(d, e, f, g, h, a, b, c, 0, _K[13]); - round(c, d, e, f, g, h, a, b, 0, _K[14]); - round(b, c, d, e, f, g, h, a, w[15], _K[15]); + roundSha256(a, b, c, d, e, f, g, h, w[0], _K[0]); + roundSha256(h, a, b, c, d, e, f, g, 0, _K[1]); + roundSha256(g, h, a, b, c, d, e, f, 0, _K[2]); + roundSha256(f, g, h, a, b, c, d, e, 0, _K[3]); + roundSha256(e, f, g, h, a, b, c, d, 0, _K[4]); + roundSha256(d, e, f, g, h, a, b, c, 0, _K[5]); + roundSha256(c, d, e, f, g, h, a, b, 0, _K[6]); + roundSha256(b, c, d, e, f, g, h, a, 0, _K[7]); + roundSha256(a, b, c, d, e, f, g, h, 0, _K[8]); + roundSha256(h, a, b, c, d, e, f, g, 0, _K[9]); + roundSha256(g, h, a, b, c, d, e, f, 0, _K[10]); + roundSha256(f, g, h, a, b, c, d, e, 0, _K[11]); + roundSha256(e, f, g, h, a, b, c, d, 0, _K[12]); + roundSha256(d, e, f, g, h, a, b, c, 0, _K[13]); + roundSha256(c, d, e, f, g, h, a, b, 0, _K[14]); + roundSha256(b, c, d, e, f, g, h, a, w[15], _K[15]); w[0] = w[0] + s0(0) + 0 + s1(0); w[1] = 0 + s0(0) + 0 + s1(w[15]); @@ -282,22 +282,22 @@ __device__ void sha256PublicKey(const unsigned int x[8], const unsigned int y[8] - round(a, b, c, d, e, f, g, h, w[0], _K[16]); - round(h, a, b, c, d, e, f, g, w[1], _K[17]); - round(g, h, a, b, c, d, e, f, w[2], _K[18]); - round(f, g, h, a, b, c, d, e, w[3], _K[19]); - round(e, f, g, h, a, b, c, d, w[4], _K[20]); - round(d, e, f, g, h, a, b, c, w[5], _K[21]); - round(c, d, e, f, g, h, a, b, w[6], _K[22]); - round(b, c, d, e, f, g, h, a, w[7], _K[23]); - round(a, b, c, d, e, f, g, h, w[8], _K[24]); - round(h, a, b, c, d, e, f, g, w[9], _K[25]); - round(g, h, a, b, c, d, e, f, w[10], _K[26]); - round(f, g, h, a, b, c, d, e, w[11], _K[27]); - round(e, f, g, h, a, b, c, d, w[12], _K[28]); - round(d, e, f, g, h, a, b, c, w[13], _K[29]); - round(c, d, e, f, g, h, a, b, w[14], _K[30]); - round(b, c, d, e, f, g, h, a, w[15], _K[31]); + roundSha256(a, b, c, d, e, f, g, h, w[0], _K[16]); + roundSha256(h, a, b, c, d, e, f, g, w[1], _K[17]); + roundSha256(g, h, a, b, c, d, e, f, w[2], _K[18]); + roundSha256(f, g, h, a, b, c, d, e, w[3], _K[19]); + roundSha256(e, f, g, h, a, b, c, d, w[4], _K[20]); + roundSha256(d, e, f, g, h, a, b, c, w[5], _K[21]); + roundSha256(c, d, e, f, g, h, a, b, w[6], _K[22]); + roundSha256(b, c, d, e, f, g, h, a, w[7], _K[23]); + roundSha256(a, b, c, d, e, f, g, h, w[8], _K[24]); + roundSha256(h, a, b, c, d, e, f, g, w[9], _K[25]); + roundSha256(g, h, a, b, c, d, e, f, w[10], _K[26]); + roundSha256(f, g, h, a, b, c, d, e, w[11], _K[27]); + roundSha256(e, f, g, h, a, b, c, d, w[12], _K[28]); + roundSha256(d, e, f, g, h, a, b, c, w[13], _K[29]); + roundSha256(c, d, e, f, g, h, a, b, w[14], _K[30]); + roundSha256(b, c, d, e, f, g, h, a, w[15], _K[31]); w[0] = w[0] + s0(w[1]) + w[9] + s1(w[14]); w[1] = w[1] + s0(w[2]) + w[10] + s1(w[15]); @@ -316,22 +316,22 @@ __device__ void sha256PublicKey(const unsigned int x[8], const unsigned int y[8] w[14] = w[14] + s0(w[15]) + w[7] + s1(w[12]); w[15] = w[15] + s0(w[0]) + w[8] + s1(w[13]); - round(a, b, c, d, e, f, g, h, w[0], _K[32]); - round(h, a, b, c, d, e, f, g, w[1], _K[33]); - round(g, h, a, b, c, d, e, f, w[2], _K[34]); - round(f, g, h, a, b, c, d, e, w[3], _K[35]); - round(e, f, g, h, a, b, c, d, w[4], _K[36]); - round(d, e, f, g, h, a, b, c, w[5], _K[37]); - round(c, d, e, f, g, h, a, b, w[6], _K[38]); - round(b, c, d, e, f, g, h, a, w[7], _K[39]); - round(a, b, c, d, e, f, g, h, w[8], _K[40]); - round(h, a, b, c, d, e, f, g, w[9], _K[41]); - round(g, h, a, b, c, d, e, f, w[10], _K[42]); - round(f, g, h, a, b, c, d, e, w[11], _K[43]); - round(e, f, g, h, a, b, c, d, w[12], _K[44]); - round(d, e, f, g, h, a, b, c, w[13], _K[45]); - round(c, d, e, f, g, h, a, b, w[14], _K[46]); - round(b, c, d, e, f, g, h, a, w[15], _K[47]); + roundSha256(a, b, c, d, e, f, g, h, w[0], _K[32]); + roundSha256(h, a, b, c, d, e, f, g, w[1], _K[33]); + roundSha256(g, h, a, b, c, d, e, f, w[2], _K[34]); + roundSha256(f, g, h, a, b, c, d, e, w[3], _K[35]); + roundSha256(e, f, g, h, a, b, c, d, w[4], _K[36]); + roundSha256(d, e, f, g, h, a, b, c, w[5], _K[37]); + roundSha256(c, d, e, f, g, h, a, b, w[6], _K[38]); + roundSha256(b, c, d, e, f, g, h, a, w[7], _K[39]); + roundSha256(a, b, c, d, e, f, g, h, w[8], _K[40]); + roundSha256(h, a, b, c, d, e, f, g, w[9], _K[41]); + roundSha256(g, h, a, b, c, d, e, f, w[10], _K[42]); + roundSha256(f, g, h, a, b, c, d, e, w[11], _K[43]); + roundSha256(e, f, g, h, a, b, c, d, w[12], _K[44]); + roundSha256(d, e, f, g, h, a, b, c, w[13], _K[45]); + roundSha256(c, d, e, f, g, h, a, b, w[14], _K[46]); + roundSha256(b, c, d, e, f, g, h, a, w[15], _K[47]); w[0] = w[0] + s0(w[1]) + w[9] + s1(w[14]); w[1] = w[1] + s0(w[2]) + w[10] + s1(w[15]); @@ -350,22 +350,22 @@ __device__ void sha256PublicKey(const unsigned int x[8], const unsigned int y[8] w[14] = w[14] + s0(w[15]) + w[7] + s1(w[12]); w[15] = w[15] + s0(w[0]) + w[8] + s1(w[13]); - round(a, b, c, d, e, f, g, h, w[0], _K[48]); - round(h, a, b, c, d, e, f, g, w[1], _K[49]); - round(g, h, a, b, c, d, e, f, w[2], _K[50]); - round(f, g, h, a, b, c, d, e, w[3], _K[51]); - round(e, f, g, h, a, b, c, d, w[4], _K[52]); - round(d, e, f, g, h, a, b, c, w[5], _K[53]); - round(c, d, e, f, g, h, a, b, w[6], _K[54]); - round(b, c, d, e, f, g, h, a, w[7], _K[55]); - round(a, b, c, d, e, f, g, h, w[8], _K[56]); - round(h, a, b, c, d, e, f, g, w[9], _K[57]); - round(g, h, a, b, c, d, e, f, w[10], _K[58]); - round(f, g, h, a, b, c, d, e, w[11], _K[59]); - round(e, f, g, h, a, b, c, d, w[12], _K[60]); - round(d, e, f, g, h, a, b, c, w[13], _K[61]); - round(c, d, e, f, g, h, a, b, w[14], _K[62]); - round(b, c, d, e, f, g, h, a, w[15], _K[63]); + roundSha256(a, b, c, d, e, f, g, h, w[0], _K[48]); + roundSha256(h, a, b, c, d, e, f, g, w[1], _K[49]); + roundSha256(g, h, a, b, c, d, e, f, w[2], _K[50]); + roundSha256(f, g, h, a, b, c, d, e, w[3], _K[51]); + roundSha256(e, f, g, h, a, b, c, d, w[4], _K[52]); + roundSha256(d, e, f, g, h, a, b, c, w[5], _K[53]); + roundSha256(c, d, e, f, g, h, a, b, w[6], _K[54]); + roundSha256(b, c, d, e, f, g, h, a, w[7], _K[55]); + roundSha256(a, b, c, d, e, f, g, h, w[8], _K[56]); + roundSha256(h, a, b, c, d, e, f, g, w[9], _K[57]); + roundSha256(g, h, a, b, c, d, e, f, w[10], _K[58]); + roundSha256(f, g, h, a, b, c, d, e, w[11], _K[59]); + roundSha256(e, f, g, h, a, b, c, d, w[12], _K[60]); + roundSha256(d, e, f, g, h, a, b, c, w[13], _K[61]); + roundSha256(c, d, e, f, g, h, a, b, w[14], _K[62]); + roundSha256(b, c, d, e, f, g, h, a, w[15], _K[63]); digest[0] = tmp[0] + a; digest[1] = tmp[1] + b; @@ -404,22 +404,22 @@ __device__ void sha256PublicKeyCompressed(const unsigned int x[8], unsigned int g = _IV[6]; h = _IV[7]; - round(a, b, c, d, e, f, g, h, w[0], _K[0]); - round(h, a, b, c, d, e, f, g, w[1], _K[1]); - round(g, h, a, b, c, d, e, f, w[2], _K[2]); - round(f, g, h, a, b, c, d, e, w[3], _K[3]); - round(e, f, g, h, a, b, c, d, w[4], _K[4]); - round(d, e, f, g, h, a, b, c, w[5], _K[5]); - round(c, d, e, f, g, h, a, b, w[6], _K[6]); - round(b, c, d, e, f, g, h, a, w[7], _K[7]); - round(a, b, c, d, e, f, g, h, w[8], _K[8]); - round(h, a, b, c, d, e, f, g, 0, _K[9]); - round(g, h, a, b, c, d, e, f, 0, _K[10]); - round(f, g, h, a, b, c, d, e, 0, _K[11]); - round(e, f, g, h, a, b, c, d, 0, _K[12]); - round(d, e, f, g, h, a, b, c, 0, _K[13]); - round(c, d, e, f, g, h, a, b, 0, _K[14]); - round(b, c, d, e, f, g, h, a, w[15], _K[15]); + roundSha256(a, b, c, d, e, f, g, h, w[0], _K[0]); + roundSha256(h, a, b, c, d, e, f, g, w[1], _K[1]); + roundSha256(g, h, a, b, c, d, e, f, w[2], _K[2]); + roundSha256(f, g, h, a, b, c, d, e, w[3], _K[3]); + roundSha256(e, f, g, h, a, b, c, d, w[4], _K[4]); + roundSha256(d, e, f, g, h, a, b, c, w[5], _K[5]); + roundSha256(c, d, e, f, g, h, a, b, w[6], _K[6]); + roundSha256(b, c, d, e, f, g, h, a, w[7], _K[7]); + roundSha256(a, b, c, d, e, f, g, h, w[8], _K[8]); + roundSha256(h, a, b, c, d, e, f, g, 0, _K[9]); + roundSha256(g, h, a, b, c, d, e, f, 0, _K[10]); + roundSha256(f, g, h, a, b, c, d, e, 0, _K[11]); + roundSha256(e, f, g, h, a, b, c, d, 0, _K[12]); + roundSha256(d, e, f, g, h, a, b, c, 0, _K[13]); + roundSha256(c, d, e, f, g, h, a, b, 0, _K[14]); + roundSha256(b, c, d, e, f, g, h, a, w[15], _K[15]); w[0] = w[0] + s0(w[1]) + 0 + s1(0); w[1] = w[1] + s0(w[2]) + 0 + s1(w[15]); @@ -438,22 +438,22 @@ __device__ void sha256PublicKeyCompressed(const unsigned int x[8], unsigned int w[14] = 0 + s0(w[15]) + w[7] + s1(w[12]); w[15] = w[15] + s0(w[0]) + w[8] + s1(w[13]); - round(a, b, c, d, e, f, g, h, w[0], _K[16]); - round(h, a, b, c, d, e, f, g, w[1], _K[17]); - round(g, h, a, b, c, d, e, f, w[2], _K[18]); - round(f, g, h, a, b, c, d, e, w[3], _K[19]); - round(e, f, g, h, a, b, c, d, w[4], _K[20]); - round(d, e, f, g, h, a, b, c, w[5], _K[21]); - round(c, d, e, f, g, h, a, b, w[6], _K[22]); - round(b, c, d, e, f, g, h, a, w[7], _K[23]); - round(a, b, c, d, e, f, g, h, w[8], _K[24]); - round(h, a, b, c, d, e, f, g, w[9], _K[25]); - round(g, h, a, b, c, d, e, f, w[10], _K[26]); - round(f, g, h, a, b, c, d, e, w[11], _K[27]); - round(e, f, g, h, a, b, c, d, w[12], _K[28]); - round(d, e, f, g, h, a, b, c, w[13], _K[29]); - round(c, d, e, f, g, h, a, b, w[14], _K[30]); - round(b, c, d, e, f, g, h, a, w[15], _K[31]); + roundSha256(a, b, c, d, e, f, g, h, w[0], _K[16]); + roundSha256(h, a, b, c, d, e, f, g, w[1], _K[17]); + roundSha256(g, h, a, b, c, d, e, f, w[2], _K[18]); + roundSha256(f, g, h, a, b, c, d, e, w[3], _K[19]); + roundSha256(e, f, g, h, a, b, c, d, w[4], _K[20]); + roundSha256(d, e, f, g, h, a, b, c, w[5], _K[21]); + roundSha256(c, d, e, f, g, h, a, b, w[6], _K[22]); + roundSha256(b, c, d, e, f, g, h, a, w[7], _K[23]); + roundSha256(a, b, c, d, e, f, g, h, w[8], _K[24]); + roundSha256(h, a, b, c, d, e, f, g, w[9], _K[25]); + roundSha256(g, h, a, b, c, d, e, f, w[10], _K[26]); + roundSha256(f, g, h, a, b, c, d, e, w[11], _K[27]); + roundSha256(e, f, g, h, a, b, c, d, w[12], _K[28]); + roundSha256(d, e, f, g, h, a, b, c, w[13], _K[29]); + roundSha256(c, d, e, f, g, h, a, b, w[14], _K[30]); + roundSha256(b, c, d, e, f, g, h, a, w[15], _K[31]); w[0] = w[0] + s0(w[1]) + w[9] + s1(w[14]); w[1] = w[1] + s0(w[2]) + w[10] + s1(w[15]); @@ -472,22 +472,22 @@ __device__ void sha256PublicKeyCompressed(const unsigned int x[8], unsigned int w[14] = w[14] + s0(w[15]) + w[7] + s1(w[12]); w[15] = w[15] + s0(w[0]) + w[8] + s1(w[13]); - round(a, b, c, d, e, f, g, h, w[0], _K[32]); - round(h, a, b, c, d, e, f, g, w[1], _K[33]); - round(g, h, a, b, c, d, e, f, w[2], _K[34]); - round(f, g, h, a, b, c, d, e, w[3], _K[35]); - round(e, f, g, h, a, b, c, d, w[4], _K[36]); - round(d, e, f, g, h, a, b, c, w[5], _K[37]); - round(c, d, e, f, g, h, a, b, w[6], _K[38]); - round(b, c, d, e, f, g, h, a, w[7], _K[39]); - round(a, b, c, d, e, f, g, h, w[8], _K[40]); - round(h, a, b, c, d, e, f, g, w[9], _K[41]); - round(g, h, a, b, c, d, e, f, w[10], _K[42]); - round(f, g, h, a, b, c, d, e, w[11], _K[43]); - round(e, f, g, h, a, b, c, d, w[12], _K[44]); - round(d, e, f, g, h, a, b, c, w[13], _K[45]); - round(c, d, e, f, g, h, a, b, w[14], _K[46]); - round(b, c, d, e, f, g, h, a, w[15], _K[47]); + roundSha256(a, b, c, d, e, f, g, h, w[0], _K[32]); + roundSha256(h, a, b, c, d, e, f, g, w[1], _K[33]); + roundSha256(g, h, a, b, c, d, e, f, w[2], _K[34]); + roundSha256(f, g, h, a, b, c, d, e, w[3], _K[35]); + roundSha256(e, f, g, h, a, b, c, d, w[4], _K[36]); + roundSha256(d, e, f, g, h, a, b, c, w[5], _K[37]); + roundSha256(c, d, e, f, g, h, a, b, w[6], _K[38]); + roundSha256(b, c, d, e, f, g, h, a, w[7], _K[39]); + roundSha256(a, b, c, d, e, f, g, h, w[8], _K[40]); + roundSha256(h, a, b, c, d, e, f, g, w[9], _K[41]); + roundSha256(g, h, a, b, c, d, e, f, w[10], _K[42]); + roundSha256(f, g, h, a, b, c, d, e, w[11], _K[43]); + roundSha256(e, f, g, h, a, b, c, d, w[12], _K[44]); + roundSha256(d, e, f, g, h, a, b, c, w[13], _K[45]); + roundSha256(c, d, e, f, g, h, a, b, w[14], _K[46]); + roundSha256(b, c, d, e, f, g, h, a, w[15], _K[47]); w[0] = w[0] + s0(w[1]) + w[9] + s1(w[14]); @@ -507,22 +507,22 @@ __device__ void sha256PublicKeyCompressed(const unsigned int x[8], unsigned int w[14] = w[14] + s0(w[15]) + w[7] + s1(w[12]); w[15] = w[15] + s0(w[0]) + w[8] + s1(w[13]); - round(a, b, c, d, e, f, g, h, w[0], _K[48]); - round(h, a, b, c, d, e, f, g, w[1], _K[49]); - round(g, h, a, b, c, d, e, f, w[2], _K[50]); - round(f, g, h, a, b, c, d, e, w[3], _K[51]); - round(e, f, g, h, a, b, c, d, w[4], _K[52]); - round(d, e, f, g, h, a, b, c, w[5], _K[53]); - round(c, d, e, f, g, h, a, b, w[6], _K[54]); - round(b, c, d, e, f, g, h, a, w[7], _K[55]); - round(a, b, c, d, e, f, g, h, w[8], _K[56]); - round(h, a, b, c, d, e, f, g, w[9], _K[57]); - round(g, h, a, b, c, d, e, f, w[10], _K[58]); - round(f, g, h, a, b, c, d, e, w[11], _K[59]); - round(e, f, g, h, a, b, c, d, w[12], _K[60]); - round(d, e, f, g, h, a, b, c, w[13], _K[61]); - round(c, d, e, f, g, h, a, b, w[14], _K[62]); - round(b, c, d, e, f, g, h, a, w[15], _K[63]); + roundSha256(a, b, c, d, e, f, g, h, w[0], _K[48]); + roundSha256(h, a, b, c, d, e, f, g, w[1], _K[49]); + roundSha256(g, h, a, b, c, d, e, f, w[2], _K[50]); + roundSha256(f, g, h, a, b, c, d, e, w[3], _K[51]); + roundSha256(e, f, g, h, a, b, c, d, w[4], _K[52]); + roundSha256(d, e, f, g, h, a, b, c, w[5], _K[53]); + roundSha256(c, d, e, f, g, h, a, b, w[6], _K[54]); + roundSha256(b, c, d, e, f, g, h, a, w[7], _K[55]); + roundSha256(a, b, c, d, e, f, g, h, w[8], _K[56]); + roundSha256(h, a, b, c, d, e, f, g, w[9], _K[57]); + roundSha256(g, h, a, b, c, d, e, f, w[10], _K[58]); + roundSha256(f, g, h, a, b, c, d, e, w[11], _K[59]); + roundSha256(e, f, g, h, a, b, c, d, w[12], _K[60]); + roundSha256(d, e, f, g, h, a, b, c, w[13], _K[61]); + roundSha256(c, d, e, f, g, h, a, b, w[14], _K[62]); + roundSha256(b, c, d, e, f, g, h, a, w[15], _K[63]); a += _IV[0]; b += _IV[1]; @@ -542,4 +542,4 @@ __device__ void sha256PublicKeyCompressed(const unsigned int x[8], unsigned int digest[6] = g; digest[7] = h; } -#endif \ No newline at end of file +#endif From 78a2dd8ccc278ee3008c411ec2ff984c2d14fd70 Mon Sep 17 00:00:00 2001 From: ByLamacq <52816027+ByLamacq@users.noreply.github.com> Date: Mon, 26 Apr 2021 09:52:37 +0200 Subject: [PATCH 05/10] Little update sha256.cuh : unsigned e -> unsigned int e in roundSha256() --- cudaMath/sha256.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cudaMath/sha256.cuh b/cudaMath/sha256.cuh index f3a956d7..2eae5c84 100644 --- a/cudaMath/sha256.cuh +++ b/cudaMath/sha256.cuh @@ -56,7 +56,7 @@ __device__ __forceinline__ unsigned int s1(unsigned int x) } -__device__ __forceinline__ void roundSha256(unsigned int a, unsigned int b, unsigned int c, unsigned int &d, unsigned e, unsigned int f, unsigned int g, unsigned int &h, unsigned int m, unsigned int k) +__device__ __forceinline__ void roundSha256(unsigned int a, unsigned int b, unsigned int c, unsigned int &d, unsigned int e, unsigned int f, unsigned int g, unsigned int &h, unsigned int m, unsigned int k) { unsigned int s = CH(e, f, g) + (rotr(e, 6) ^ rotr(e, 11) ^ rotr(e, 25)) + k + m; From 23258556c9e43cce6c9c737974ed24714e3c2de7 Mon Sep 17 00:00:00 2001 From: blm Date: Fri, 7 May 2021 00:21:26 +0200 Subject: [PATCH 06/10] Update to nvidia drivers 460+ and cuda 11 for linux. Many changes, see changelog file for more informations. --- CudaKeySearchDevice/CudaDeviceKeys.cu | 20 ++- CudaKeySearchDevice/CudaDeviceKeys.h | 5 +- CudaKeySearchDevice/CudaKeyCheckDevice.cu | 167 ++++++++++++++++++ CudaKeySearchDevice/CudaKeySearchDevice.cpp | 8 +- CudaKeySearchDevice/CudaKeySearchDevice.cu | 177 +------------------- CudaKeySearchDevice/cudabridge.cu | 35 ++-- CudaKeySearchDevice/cudabridge.h | 5 +- cudaMath/sha256.cuh | 4 +- 8 files changed, 230 insertions(+), 191 deletions(-) create mode 100644 CudaKeySearchDevice/CudaKeyCheckDevice.cu diff --git a/CudaKeySearchDevice/CudaDeviceKeys.cu b/CudaKeySearchDevice/CudaDeviceKeys.cu index 85db71b5..e4183ad6 100644 --- a/CudaKeySearchDevice/CudaDeviceKeys.cu +++ b/CudaKeySearchDevice/CudaDeviceKeys.cu @@ -24,6 +24,20 @@ __device__ unsigned int *ec::getYPtr() __global__ void multiplyStepKernel(unsigned int *privateKeys, int pointsPerThread, int step, unsigned int *chain, const unsigned int *gxPtr, const unsigned int *gyPtr); +unsigned int * CudaDeviceKeys::getXPoints() +{ + return _devX; +} + +unsigned int * CudaDeviceKeys::getYPoints() +{ + return _devY; +} + +unsigned int * CudaDeviceKeys::getChain() +{ + return _devChain; +} int CudaDeviceKeys::getIndex(int block, int thread, int idx) { @@ -269,9 +283,11 @@ cudaError_t CudaDeviceKeys::init(int blocks, int threads, int pointsPerThread, c void CudaDeviceKeys::clearPublicKeys() { + cudaFree(_devChain); cudaFree(_devX); cudaFree(_devY); + _devChain = NULL; _devX = NULL; _devY = NULL; } @@ -281,9 +297,9 @@ void CudaDeviceKeys::clearPrivateKeys() cudaFree(_devBasePointX); cudaFree(_devBasePointY); cudaFree(_devPrivate); - cudaFree(_devChain); + //cudaFree(_devChain); - _devChain = NULL; + //_devChain = NULL; _devBasePointX = NULL; _devBasePointY = NULL; _devPrivate = NULL; diff --git a/CudaKeySearchDevice/CudaDeviceKeys.h b/CudaKeySearchDevice/CudaDeviceKeys.h index f2407f10..8dba2f6d 100644 --- a/CudaKeySearchDevice/CudaDeviceKeys.h +++ b/CudaKeySearchDevice/CudaDeviceKeys.h @@ -76,6 +76,9 @@ class CudaDeviceKeys { void clearPublicKeys(); + unsigned int * getXPoints(); + unsigned int * getYPoints(); + unsigned int * getChain(); }; -#endif \ No newline at end of file +#endif diff --git a/CudaKeySearchDevice/CudaKeyCheckDevice.cu b/CudaKeySearchDevice/CudaKeyCheckDevice.cu new file mode 100644 index 00000000..082085c4 --- /dev/null +++ b/CudaKeySearchDevice/CudaKeyCheckDevice.cu @@ -0,0 +1,167 @@ +#include +#include +#include +#include "KeySearchTypes.h" +#include "CudaKeySearchDevice.h" +#include "ptx.cuh" +#include "secp256k1.cuh" + +#include "sha256.cuh" +#include "ripemd160.cuh" + +#include "secp256k1.h" + +#include "CudaHashLookup.cuh" +#include "CudaAtomicList.cuh" +#include "CudaDeviceKeys.cuh" + + +__device__ void doRMD160FinalRound(const unsigned int hIn[5], unsigned int hOut[5]) +{ + const unsigned int iv[5] = { + 0x67452301, + 0xefcdab89, + 0x98badcfe, + 0x10325476, + 0xc3d2e1f0 + }; + + for(int i = 0; i < 5; i++) { + hOut[i] = endian(hIn[i] + iv[(i + 1) % 5]); + } +} + +__device__ void hashPublicKey(const unsigned int *x, const unsigned int *y, unsigned int *digestOut) +{ + unsigned int hash[8]; + + sha256PublicKey(x, y, hash); + + // Swap to little-endian + for(int i = 0; i < 8; i++) { + hash[i] = endian(hash[i]); + } + + ripemd160sha256NoFinal(hash, digestOut); +} + +__device__ void hashPublicKeyCompressed(const unsigned int *x, unsigned int yParity, unsigned int *digestOut) +{ + unsigned int hash[8]; + + sha256PublicKeyCompressed(x, yParity, hash); + + // Swap to little-endian + for(int i = 0; i < 8; i++) { + hash[i] = endian(hash[i]); + } + + ripemd160sha256NoFinal(hash, digestOut); +} + +__device__ void setResultFound(const int numBlocks, bool compressed, unsigned int x[8], unsigned int y[8], unsigned int digest[5]) +{ + CudaDeviceResult r; + + unsigned int virtualPointsPerThread = blockIdx.x / numBlocks; // ex: (uint) 8 / 6 = 1 + unsigned int virtualBlock = blockIdx.x % numBlocks; + + r.block = virtualBlock; + r.thread = threadIdx.x; + r.idx = virtualPointsPerThread; + r.compressed = compressed; + + for(int i = 0; i < 8; i++) { + r.x[i] = x[i]; + r.y[i] = y[i]; + } + + doRMD160FinalRound(digest, r.digest); + + atomicListAdd(&r, sizeof(r)); +} + +/** + * Reads an 2-vector4 big integer from device memory with virtualized loop + */ +__device__ static void readIntVirtualized(unsigned int *ara, const int numBlocks, unsigned int x[8]) +{ + uint4 *araTmp = reinterpret_cast(ara); + + unsigned int virtualPointsPerThread = blockIdx.x / numBlocks; // ex: (uint) 8 / 6 = 1 + unsigned int virtualBlock = blockIdx.x % numBlocks; // ex: 8 % 6 = 2 + unsigned int totalThreads = numBlocks * blockDim.x; // 6 * 128 = 768 + unsigned int base = virtualPointsPerThread * totalThreads * 2; // 1 * 768 + unsigned int threadId = blockDim.x * virtualBlock + threadIdx.x; // 128 * 2 * ... + unsigned int index = base + threadId; + + uint4 xTmp = araTmp[index]; + x[0] = xTmp.x; + x[1] = xTmp.y; + x[2] = xTmp.z; + x[3] = xTmp.w; + + index += totalThreads; + + xTmp = araTmp[index]; + x[4] = xTmp.x; + x[5] = xTmp.y; + x[6] = xTmp.z; + x[7] = xTmp.w; +} + +__device__ static unsigned int readIntLSWVirtualized(unsigned int *ara, const int numBlocks) +{ + uint4 *araTmp = reinterpret_cast(ara); + + unsigned int virtualPointsPerThread = blockIdx.x / numBlocks; // ex: (uint) 8 / 6 = 1 + unsigned int virtualBlock = blockIdx.x % numBlocks; // ex: 8 % 6 = 2 + unsigned int totalThreads = numBlocks * blockDim.x; // 6 * 128 = 768 + unsigned int base = virtualPointsPerThread * totalThreads * 2; // 1 * 768 + unsigned int threadId = blockDim.x * virtualBlock + threadIdx.x; // 128 * 2 * ... + unsigned int index = base + threadId; + + index += totalThreads; + + uint4 xTmp = araTmp[index]; + + return xTmp.w; +} + +__device__ void hashAndCheck(unsigned int *xPtr, unsigned int *yPtr, const int numBlocks, const int compression) +{ + unsigned int x[8]; + unsigned int digest[5]; + + readIntVirtualized(xPtr, numBlocks, x); + + if(compression == PointCompressionType::UNCOMPRESSED || compression == PointCompressionType::BOTH) { + unsigned int y[8]; + readIntVirtualized(yPtr, numBlocks, y); + + hashPublicKey(x, y, digest); + + if(checkHash(digest)) { + setResultFound(numBlocks, false, x, y, digest); + } + } + + if(compression == PointCompressionType::COMPRESSED || compression == PointCompressionType::BOTH) { + hashPublicKeyCompressed(x, readIntLSWVirtualized(yPtr, numBlocks), digest); + + if(checkHash(digest)) { + unsigned int y[8]; + readIntVirtualized(yPtr, numBlocks, y); + setResultFound(numBlocks, true, x, y, digest); + } + } +} + +/** +* Performs a single iteration +*/ +__global__ void keyCheckKernel(unsigned int *xPtr, unsigned int *yPtr, const int blocks, const int compression) +{ + hashAndCheck(xPtr, yPtr, blocks, compression); +} + diff --git a/CudaKeySearchDevice/CudaKeySearchDevice.cpp b/CudaKeySearchDevice/CudaKeySearchDevice.cpp index aad1fd31..c9569ff6 100644 --- a/CudaKeySearchDevice/CudaKeySearchDevice.cpp +++ b/CudaKeySearchDevice/CudaKeySearchDevice.cpp @@ -80,7 +80,7 @@ void CudaKeySearchDevice::init(const secp256k1::uint256 &start, int compression, generateStartingPoints(); - cudaCall(allocateChainBuf(_threads * _blocks * _pointsPerThread)); + //cudaCall(allocateChainBuf(_threads * _blocks * _pointsPerThread)); // Set the incrementor secp256k1::ecpoint g = secp256k1::G(); @@ -148,9 +148,9 @@ void CudaKeySearchDevice::doStep() try { if(_iterations < 2 && _startExponent.cmp(numKeys) <= 0) { - callKeyFinderKernel(_blocks, _threads, _pointsPerThread, true, _compression); + callKeyFinderKernel(_blocks, _threads, _pointsPerThread, _deviceKeys.getXPoints(), _deviceKeys.getYPoints(), _deviceKeys.getChain(), true, _compression); } else { - callKeyFinderKernel(_blocks, _threads, _pointsPerThread, false, _compression); + callKeyFinderKernel(_blocks, _threads, _pointsPerThread, _deviceKeys.getXPoints(), _deviceKeys.getYPoints(), _deviceKeys.getChain(), false, _compression); } } catch(cuda::CudaException ex) { throw KeySearchException(ex.msg); @@ -313,4 +313,4 @@ secp256k1::uint256 CudaKeySearchDevice::getNextKey() uint64_t totalPoints = (uint64_t)_pointsPerThread * _threads * _blocks; return _startExponent + secp256k1::uint256(totalPoints) * _iterations * _stride; -} \ No newline at end of file +} diff --git a/CudaKeySearchDevice/CudaKeySearchDevice.cu b/CudaKeySearchDevice/CudaKeySearchDevice.cu index cbb79ad6..30d3ebac 100644 --- a/CudaKeySearchDevice/CudaKeySearchDevice.cu +++ b/CudaKeySearchDevice/CudaKeySearchDevice.cu @@ -3,12 +3,10 @@ #include #include "KeySearchTypes.h" #include "CudaKeySearchDevice.h" + #include "ptx.cuh" #include "secp256k1.cuh" -#include "sha256.cuh" -#include "ripemd160.cuh" - #include "secp256k1.h" #include "CudaHashLookup.cuh" @@ -16,57 +14,8 @@ #include "CudaDeviceKeys.cuh" __constant__ unsigned int _INC_X[8]; - __constant__ unsigned int _INC_Y[8]; -__constant__ unsigned int *_CHAIN[1]; - -static unsigned int *_chainBufferPtr = NULL; - - -__device__ void doRMD160FinalRound(const unsigned int hIn[5], unsigned int hOut[5]) -{ - const unsigned int iv[5] = { - 0x67452301, - 0xefcdab89, - 0x98badcfe, - 0x10325476, - 0xc3d2e1f0 - }; - - for(int i = 0; i < 5; i++) { - hOut[i] = endian(hIn[i] + iv[(i + 1) % 5]); - } -} - - -/** - * Allocates device memory for storing the multiplication chain used in - the batch inversion operation - */ -cudaError_t allocateChainBuf(unsigned int count) -{ - cudaError_t err = cudaMalloc(&_chainBufferPtr, count * sizeof(unsigned int) * 8); - - if(err) { - return err; - } - - err = cudaMemcpyToSymbol(_CHAIN, &_chainBufferPtr, sizeof(unsigned int *)); - if(err) { - cudaFree(_chainBufferPtr); - } - - return err; -} - -void cleanupChainBuf() -{ - if(_chainBufferPtr != NULL) { - cudaFree(_chainBufferPtr); - _chainBufferPtr = NULL; - } -} /** *Sets the EC point which all points will be incremented by @@ -87,92 +36,14 @@ cudaError_t setIncrementorPoint(const secp256k1::uint256 &x, const secp256k1::ui return cudaMemcpyToSymbol(_INC_Y, yWords, sizeof(unsigned int) * 8); } - - -__device__ void hashPublicKey(const unsigned int *x, const unsigned int *y, unsigned int *digestOut) -{ - unsigned int hash[8]; - - sha256PublicKey(x, y, hash); - - // Swap to little-endian - for(int i = 0; i < 8; i++) { - hash[i] = endian(hash[i]); - } - - ripemd160sha256NoFinal(hash, digestOut); -} - -__device__ void hashPublicKeyCompressed(const unsigned int *x, unsigned int yParity, unsigned int *digestOut) -{ - unsigned int hash[8]; - - sha256PublicKeyCompressed(x, yParity, hash); - - // Swap to little-endian - for(int i = 0; i < 8; i++) { - hash[i] = endian(hash[i]); - } - - ripemd160sha256NoFinal(hash, digestOut); -} - - -__device__ void setResultFound(int idx, bool compressed, unsigned int x[8], unsigned int y[8], unsigned int digest[5]) -{ - CudaDeviceResult r; - - r.block = blockIdx.x; - r.thread = threadIdx.x; - r.idx = idx; - r.compressed = compressed; - - for(int i = 0; i < 8; i++) { - r.x[i] = x[i]; - r.y[i] = y[i]; - } - - doRMD160FinalRound(digest, r.digest); - - atomicListAdd(&r, sizeof(r)); -} - -__device__ void doIteration(int pointsPerThread, int compression) +__device__ void doIteration(unsigned int *xPtr, unsigned int *yPtr, unsigned int *chain, int pointsPerThread) { - unsigned int *chain = _CHAIN[0]; - unsigned int *xPtr = ec::getXPtr(); - unsigned int *yPtr = ec::getYPtr(); - // Multiply together all (_Gx - x) and then invert unsigned int inverse[8] = {0,0,0,0,0,0,0,1}; for(int i = 0; i < pointsPerThread; i++) { unsigned int x[8]; - - unsigned int digest[5]; - readInt(xPtr, i, x); - if(compression == PointCompressionType::UNCOMPRESSED || compression == PointCompressionType::BOTH) { - unsigned int y[8]; - readInt(yPtr, i, y); - - hashPublicKey(x, y, digest); - - if(checkHash(digest)) { - setResultFound(i, false, x, y, digest); - } - } - - if(compression == PointCompressionType::COMPRESSED || compression == PointCompressionType::BOTH) { - hashPublicKeyCompressed(x, readIntLSW(yPtr, i), digest); - - if(checkHash(digest)) { - unsigned int y[8]; - readInt(yPtr, i, y); - setResultFound(i, true, x, y, digest); - } - } - beginBatchAdd(_INC_X, x, chain, i, i, inverse); } @@ -190,46 +61,14 @@ __device__ void doIteration(int pointsPerThread, int compression) } } -__device__ void doIterationWithDouble(int pointsPerThread, int compression) +__device__ void doIterationWithDouble(unsigned int *xPtr, unsigned int *yPtr, unsigned int *chain, int pointsPerThread) { - unsigned int *chain = _CHAIN[0]; - unsigned int *xPtr = ec::getXPtr(); - unsigned int *yPtr = ec::getYPtr(); - // Multiply together all (_Gx - x) and then invert unsigned int inverse[8] = {0,0,0,0,0,0,0,1}; for(int i = 0; i < pointsPerThread; i++) { unsigned int x[8]; - - unsigned int digest[5]; - readInt(xPtr, i, x); - // uncompressed - if(compression == PointCompressionType::UNCOMPRESSED || compression == PointCompressionType::BOTH) { - unsigned int y[8]; - readInt(yPtr, i, y); - hashPublicKey(x, y, digest); - - if(checkHash(digest)) { - setResultFound(i, false, x, y, digest); - } - } - - // compressed - if(compression == PointCompressionType::COMPRESSED || compression == PointCompressionType::BOTH) { - - hashPublicKeyCompressed(x, readIntLSW(yPtr, i), digest); - - if(checkHash(digest)) { - - unsigned int y[8]; - readInt(yPtr, i, y); - - setResultFound(i, true, x, y, digest); - } - } - beginBatchAddWithDouble(_INC_X, _INC_Y, xPtr, chain, i, i, inverse); } @@ -250,12 +89,12 @@ __device__ void doIterationWithDouble(int pointsPerThread, int compression) /** * Performs a single iteration */ -__global__ void keyFinderKernel(int points, int compression) +__global__ void keyFinderKernel(unsigned int *xPtr, unsigned int *yPtr, unsigned int *chainPtr, int points) { - doIteration(points, compression); + doIteration(xPtr, yPtr, chainPtr, points); } -__global__ void keyFinderKernelWithDouble(int points, int compression) +__global__ void keyFinderKernelWithDouble(unsigned int *xPtr, unsigned int *yPtr, unsigned int *chainPtr, int points) { - doIterationWithDouble(points, compression); -} \ No newline at end of file + doIterationWithDouble(xPtr, yPtr, chainPtr, points); +} diff --git a/CudaKeySearchDevice/cudabridge.cu b/CudaKeySearchDevice/cudabridge.cu index 33325d97..63fa7079 100644 --- a/CudaKeySearchDevice/cudabridge.cu +++ b/CudaKeySearchDevice/cudabridge.cu @@ -1,19 +1,34 @@ #include "cudabridge.h" -__global__ void keyFinderKernel(int points, int compression); -__global__ void keyFinderKernelWithDouble(int points, int compression); +__global__ void keyCheckKernel(unsigned int *xPtr, unsigned int *yPtr, int blocks, int compression); +__global__ void keyFinderKernel(unsigned int *xPtr, unsigned int *yPtr, unsigned int *chainPtr, int points); +__global__ void keyFinderKernelWithDouble(unsigned int *xPtr, unsigned int *yPtr, unsigned int *chainPtr, int points); -void callKeyFinderKernel(int blocks, int threads, int points, bool useDouble, int compression) +void callKeyFinderKernel(int blocks, int threads, int points, unsigned int *xPtr, unsigned int *yPtr, unsigned int *chainPtr, bool useDouble, int compression) { - if(useDouble) { - keyFinderKernelWithDouble <<>>(points, compression); - } else { - keyFinderKernel <<>> (points, compression); - } - waitForKernel(); + int blocksMultPoints = blocks * points; + + keyCheckKernel <<>> (xPtr, yPtr, blocks, compression); + checkKernelLaunch(); + + if(useDouble) { + keyFinderKernelWithDouble <<>> (xPtr, yPtr, chainPtr, points); + } else { + keyFinderKernel <<>> (xPtr, yPtr, chainPtr, points); + } + waitForKernel(); } +void checkKernelLaunch() +{ + // Check for kernel launch error + cudaError_t err = cudaGetLastError(); + + if(err != cudaSuccess) { + throw cuda::CudaException(err); + } +} void waitForKernel() { @@ -30,4 +45,4 @@ void waitForKernel() if(err != cudaSuccess) { throw cuda::CudaException(err); } -} \ No newline at end of file +} diff --git a/CudaKeySearchDevice/cudabridge.h b/CudaKeySearchDevice/cudabridge.h index eaafe3a2..3a5e51b0 100644 --- a/CudaKeySearchDevice/cudabridge.h +++ b/CudaKeySearchDevice/cudabridge.h @@ -8,12 +8,13 @@ #include "secp256k1.h" -void callKeyFinderKernel(int blocks, int threads, int points, bool useDouble, int compression); +void callKeyFinderKernel(int blocks, int threads, int points, unsigned int *xPtr, unsigned int *yPtr, unsigned int *chainPtr, bool useDouble, int compression); +void checkKernelLaunch(); void waitForKernel(); cudaError_t setIncrementorPoint(const secp256k1::uint256 &x, const secp256k1::uint256 &y); cudaError_t allocateChainBuf(unsigned int count); void cleanupChainBuf(); -#endif \ No newline at end of file +#endif diff --git a/cudaMath/sha256.cuh b/cudaMath/sha256.cuh index f3a956d7..a66b80b7 100644 --- a/cudaMath/sha256.cuh +++ b/cudaMath/sha256.cuh @@ -4,8 +4,6 @@ #include #include -#include - __constant__ unsigned int _K[64] = { 0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, 0x923f82a4, 0xab1c5ed5, @@ -56,7 +54,7 @@ __device__ __forceinline__ unsigned int s1(unsigned int x) } -__device__ __forceinline__ void roundSha256(unsigned int a, unsigned int b, unsigned int c, unsigned int &d, unsigned e, unsigned int f, unsigned int g, unsigned int &h, unsigned int m, unsigned int k) +__device__ __forceinline__ void roundSha256(unsigned int a, unsigned int b, unsigned int c, unsigned int &d, unsigned int e, unsigned int f, unsigned int g, unsigned int &h, unsigned int m, unsigned int k) { unsigned int s = CH(e, f, g) + (rotr(e, 6) ^ rotr(e, 11) ^ rotr(e, 25)) + k + m; From 7f0065a19752dbbfd79cdd0de329dfec589c63fe Mon Sep 17 00:00:00 2001 From: ByLamacq Date: Sun, 16 May 2021 22:30:54 +0200 Subject: [PATCH 07/10] Update visual studio config to cuda 11.3 --- BitCrack.props | 8 ++++---- CudaKeySearchDevice/CudaKeySearchDevice.vcxproj | 6 ++++-- cudaInfo/cudaInfo.vcxproj | 5 +++-- cudaMath/cudaMath.vcxproj | 5 +++-- 4 files changed, 14 insertions(+), 10 deletions(-) diff --git a/BitCrack.props b/BitCrack.props index 0f980a56..f6a0e5ef 100644 --- a/BitCrack.props +++ b/BitCrack.props @@ -2,10 +2,10 @@ - C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.1\include - C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.1\lib\x64 - C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.1\include - C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.1\lib\x64 + C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.3\include + C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.3\lib\x64 + C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.3\include + C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.3\lib\x64 diff --git a/CudaKeySearchDevice/CudaKeySearchDevice.vcxproj b/CudaKeySearchDevice/CudaKeySearchDevice.vcxproj index 490c53e3..544ca3b2 100644 --- a/CudaKeySearchDevice/CudaKeySearchDevice.vcxproj +++ b/CudaKeySearchDevice/CudaKeySearchDevice.vcxproj @@ -26,6 +26,7 @@ + @@ -56,7 +57,7 @@ - + @@ -87,6 +88,7 @@ 64 true %(CodeGeneration) + true @@ -114,6 +116,6 @@ - + \ No newline at end of file diff --git a/cudaInfo/cudaInfo.vcxproj b/cudaInfo/cudaInfo.vcxproj index 66365062..26f25d9e 100644 --- a/cudaInfo/cudaInfo.vcxproj +++ b/cudaInfo/cudaInfo.vcxproj @@ -39,7 +39,7 @@ - + @@ -65,6 +65,7 @@ 64 + true @@ -89,6 +90,6 @@ - + \ No newline at end of file diff --git a/cudaMath/cudaMath.vcxproj b/cudaMath/cudaMath.vcxproj index fcf40975..8248da4f 100644 --- a/cudaMath/cudaMath.vcxproj +++ b/cudaMath/cudaMath.vcxproj @@ -31,7 +31,7 @@ - + @@ -56,6 +56,7 @@ 64 + true @@ -85,6 +86,6 @@ - + \ No newline at end of file From a725df8c307a0e0d0f681334be914ff00114e32e Mon Sep 17 00:00:00 2001 From: blm Date: Mon, 17 May 2021 21:10:24 +0200 Subject: [PATCH 08/10] Small format change : tab to 4 space --- KeyFinder/main.cpp | 126 ++++++++++++++++++++++----------------------- 1 file changed, 63 insertions(+), 63 deletions(-) diff --git a/KeyFinder/main.cpp b/KeyFinder/main.cpp index 22c699df..eb1dc6e0 100644 --- a/KeyFinder/main.cpp +++ b/KeyFinder/main.cpp @@ -461,8 +461,8 @@ bool parseShare(const std::string &s, uint32_t &idx, uint32_t &total) int main(int argc, char **argv) { - bool optCompressed = false; - bool optUncompressed = false; + bool optCompressed = false; + bool optUncompressed = false; bool listDevices = false; bool optShares = false; bool optThreads = false; @@ -494,23 +494,23 @@ int main(int argc, char **argv) } // Check for arguments - if(argc == 1) { - usage(); - return 0; - } + if(argc == 1) { + usage(); + return 0; + } - CmdParse parser; - parser.add("-d", "--device", true); - parser.add("-t", "--threads", true); - parser.add("-b", "--blocks", true); - parser.add("-p", "--points", true); - parser.add("-d", "--device", true); - parser.add("-c", "--compressed", false); - parser.add("-u", "--uncompressed", false); + CmdParse parser; + parser.add("-d", "--device", true); + parser.add("-t", "--threads", true); + parser.add("-b", "--blocks", true); + parser.add("-p", "--points", true); + parser.add("-d", "--device", true); + parser.add("-c", "--compressed", false); + parser.add("-u", "--uncompressed", false); parser.add("", "--compression", true); - parser.add("-i", "--in", true); - parser.add("-o", "--out", true); + parser.add("-i", "--in", true); + parser.add("-o", "--out", true); parser.add("-f", "--follow", false); parser.add("", "--list-devices", false); parser.add("", "--keyspace", true); @@ -527,32 +527,32 @@ int main(int argc, char **argv) std::vector args = parser.getArgs(); - for(unsigned int i = 0; i < args.size(); i++) { - OptArg optArg = args[i]; - std::string opt = args[i].option; + for(unsigned int i = 0; i < args.size(); i++) { + OptArg optArg = args[i]; + std::string opt = args[i].option; - try { - if(optArg.equals("-t", "--threads")) { - _config.threads = util::parseUInt32(optArg.arg); + try { + if(optArg.equals("-t", "--threads")) { + _config.threads = util::parseUInt32(optArg.arg); optThreads = true; } else if(optArg.equals("-b", "--blocks")) { _config.blocks = util::parseUInt32(optArg.arg); optBlocks = true; - } else if(optArg.equals("-p", "--points")) { - _config.pointsPerThread = util::parseUInt32(optArg.arg); + } else if(optArg.equals("-p", "--points")) { + _config.pointsPerThread = util::parseUInt32(optArg.arg); optPoints = true; - } else if(optArg.equals("-d", "--device")) { - _config.device = util::parseUInt32(optArg.arg); - } else if(optArg.equals("-c", "--compressed")) { - optCompressed = true; + } else if(optArg.equals("-d", "--device")) { + _config.device = util::parseUInt32(optArg.arg); + } else if(optArg.equals("-c", "--compressed")) { + optCompressed = true; } else if(optArg.equals("-u", "--uncompressed")) { optUncompressed = true; } else if(optArg.equals("", "--compression")) { _config.compression = parseCompressionString(optArg.arg); - } else if(optArg.equals("-i", "--in")) { - _config.targetsFile = optArg.arg; - } else if(optArg.equals("-o", "--out")) { - _config.resultsFile = optArg.arg; + } else if(optArg.equals("-i", "--in")) { + _config.targetsFile = optArg.arg; + } else if(optArg.equals("-o", "--out")) { + _config.resultsFile = optArg.arg; } else if(optArg.equals("", "--list-devices")) { listDevices = true; } else if(optArg.equals("", "--continue")) { @@ -604,43 +604,43 @@ int main(int argc, char **argv) _config.follow = true; } - } catch(std::string err) { - Logger::log(LogLevel::Error, "Error " + opt + ": " + err); - return 1; - } - } + } catch(std::string err) { + Logger::log(LogLevel::Error, "Error " + opt + ": " + err); + return 1; + } + } if(listDevices) { printDeviceList(_devices); return 0; } - // Verify device exists - if(_config.device < 0 || _config.device >= _devices.size()) { - Logger::log(LogLevel::Error, "device " + util::format(_config.device) + " does not exist"); - return 1; - } + // Verify device exists + if(_config.device < 0 || _config.device >= _devices.size()) { + Logger::log(LogLevel::Error, "device " + util::format(_config.device) + " does not exist"); + return 1; + } - // Parse operands - std::vector ops = parser.getOperands(); + // Parse operands + std::vector ops = parser.getOperands(); // If there are no operands, then we must be reading from a file, otherwise // expect addresses on the commandline - if(ops.size() == 0) { - if(_config.targetsFile.length() == 0) { - Logger::log(LogLevel::Error, "Missing arguments"); - usage(); - return 1; - } - } else { - for(unsigned int i = 0; i < ops.size(); i++) { + if(ops.size() == 0) { + if(_config.targetsFile.length() == 0) { + Logger::log(LogLevel::Error, "Missing arguments"); + usage(); + return 1; + } + } else { + for(unsigned int i = 0; i < ops.size(); i++) { if(!Address::verifyAddress(ops[i])) { Logger::log(LogLevel::Error, "Invalid address '" + ops[i] + "'"); return 1; } - _config.targets.push_back(ops[i]); - } - } + _config.targets.push_back(ops[i]); + } + } // Calculate where to start and end in the keyspace when the --share option is used if(optShares) { @@ -663,14 +663,14 @@ int main(int argc, char **argv) _config.startKey = startPos; } - // Check option for compressed, uncompressed, or both - if(optCompressed && optUncompressed) { - _config.compression = PointCompressionType::BOTH; - } else if(optCompressed) { - _config.compression = PointCompressionType::COMPRESSED; - } else if(optUncompressed) { - _config.compression = PointCompressionType::UNCOMPRESSED; - } + // Check option for compressed, uncompressed, or both + if(optCompressed && optUncompressed) { + _config.compression = PointCompressionType::BOTH; + } else if(optCompressed) { + _config.compression = PointCompressionType::COMPRESSED; + } else if(optUncompressed) { + _config.compression = PointCompressionType::UNCOMPRESSED; + } if(_config.checkpointFile.length() > 0) { readCheckpointFile(); From d83f40373c3b548fb9a59ffde8588de9001d68d3 Mon Sep 17 00:00:00 2001 From: blm Date: Mon, 17 May 2021 21:15:54 +0200 Subject: [PATCH 09/10] Small format change : tab to 4 spases --- KeyFinder/main.cpp | 86 +++++++++++++++++++++++----------------------- 1 file changed, 43 insertions(+), 43 deletions(-) diff --git a/KeyFinder/main.cpp b/KeyFinder/main.cpp index eb1dc6e0..83564320 100644 --- a/KeyFinder/main.cpp +++ b/KeyFinder/main.cpp @@ -72,35 +72,35 @@ static uint64_t _startTime = 0; */ void resultCallback(KeySearchResult info) { - if(_config.resultsFile.length() != 0) { - Logger::log(LogLevel::Info, "Found key for address '" + info.address + "'. Written to '" + _config.resultsFile + "'"); + if(_config.resultsFile.length() != 0) { + Logger::log(LogLevel::Info, "Found key for address '" + info.address + "'. Written to '" + _config.resultsFile + "'"); - std::string s = info.address + " " + info.privateKey.toString(16) + " " + info.publicKey.toString(info.compressed); - util::appendToFile(_config.resultsFile, s); + std::string s = info.address + " " + info.privateKey.toString(16) + " " + info.publicKey.toString(info.compressed); + util::appendToFile(_config.resultsFile, s); - return; - } + return; + } - std::string logStr = "Address: " + info.address + "\n"; - logStr += "Private key: " + info.privateKey.toString(16) + "\n"; - logStr += "Compressed: "; + std::string logStr = "Address: " + info.address + "\n"; + logStr += "Private key: " + info.privateKey.toString(16) + "\n"; + logStr += "Compressed: "; - if(info.compressed) { - logStr += "yes\n"; - } else { - logStr += "no\n"; - } + if(info.compressed) { + logStr += "yes\n"; + } else { + logStr += "no\n"; + } - logStr += "Public key: \n"; + logStr += "Public key: \n"; - if(info.compressed) { - logStr += info.publicKey.toString(true) + "\n"; - } else { - logStr += info.publicKey.x.toString(16) + "\n"; - logStr += info.publicKey.y.toString(16) + "\n"; - } + if(info.compressed) { + logStr += info.publicKey.toString(true) + "\n"; + } else { + logStr += info.publicKey.x.toString(16) + "\n"; + logStr += info.publicKey.y.toString(16) + "\n"; + } - Logger::log(LogLevel::Info, logStr); + Logger::log(LogLevel::Info, logStr); } /** @@ -108,28 +108,28 @@ Callback to display progress */ void statusCallback(KeySearchStatus info) { - std::string speedStr; + std::string speedStr; - if(info.speed < 0.01) { - speedStr = "< 0.01 MKey/s"; - } else { - speedStr = util::format("%.2f", info.speed) + " MKey/s"; - } + if(info.speed < 0.01) { + speedStr = "< 0.01 MKey/s"; + } else { + speedStr = util::format("%.2f", info.speed) + " MKey/s"; + } - std::string totalStr = "(" + util::formatThousands(_config.totalkeys + info.total) + " total)"; + std::string totalStr = "(" + util::formatThousands(_config.totalkeys + info.total) + " total)"; - std::string timeStr = "[" + util::formatSeconds((unsigned int)((_config.elapsed + info.totalTime) / 1000)) + "]"; + std::string timeStr = "[" + util::formatSeconds((unsigned int)((_config.elapsed + info.totalTime) / 1000)) + "]"; - std::string usedMemStr = util::format((info.deviceMemory - info.freeMemory) /(1024 * 1024)); + std::string usedMemStr = util::format((info.deviceMemory - info.freeMemory) /(1024 * 1024)); - std::string totalMemStr = util::format(info.deviceMemory / (1024 * 1024)); + std::string totalMemStr = util::format(info.deviceMemory / (1024 * 1024)); std::string targetStr = util::format(info.targets) + " target" + (info.targets > 1 ? "s" : ""); - // Fit device name in 16 characters, pad with spaces if less - std::string devName = info.deviceName.substr(0, 16); - devName += std::string(16 - devName.length(), ' '); + // Fit device name in 16 characters, pad with spaces if less + std::string devName = info.deviceName.substr(0, 16); + devName += std::string(16 - devName.length(), ' '); const char *formatStr = NULL; @@ -139,7 +139,7 @@ void statusCallback(KeySearchStatus info) formatStr = "\r%s %s / %sMB | %s %s %s %s"; } - printf(formatStr, devName.c_str(), usedMemStr.c_str(), totalMemStr.c_str(), targetStr.c_str(), speedStr.c_str(), totalStr.c_str(), timeStr.c_str()); + printf(formatStr, devName.c_str(), usedMemStr.c_str(), totalMemStr.c_str(), targetStr.c_str(), speedStr.c_str(), totalStr.c_str(), timeStr.c_str()); if(_config.checkpointFile.length() > 0) { uint64_t t = util::getSystemTime(); @@ -222,19 +222,19 @@ void usage() Finds default parameters depending on the device */ typedef struct { - int threads; - int blocks; - int pointsPerThread; + int threads; + int blocks; + int pointsPerThread; }DeviceParameters; DeviceParameters getDefaultParameters(const DeviceManager::DeviceInfo &device) { - DeviceParameters p; - p.threads = 256; + DeviceParameters p; + p.threads = 256; p.blocks = 32; - p.pointsPerThread = 32; + p.pointsPerThread = 32; - return p; + return p; } static KeySearchDevice *getDeviceContext(DeviceManager::DeviceInfo &device, int blocks, int threads, int pointsPerThread) From 4860b1021ca60bb7968cb42a393387f5391024f3 Mon Sep 17 00:00:00 2001 From: blm Date: Mon, 17 May 2021 21:53:11 +0200 Subject: [PATCH 10/10] Small format change : tab to 4 spaces --- KeyFinderLib/KeyFinder.cpp | 198 ++++++++++--------- KeyFinderLib/KeyFinder.h | 48 ++--- KeyFinderLib/KeyFinderShared.h | 24 +-- util/util.cpp | 348 ++++++++++++++++----------------- 4 files changed, 308 insertions(+), 310 deletions(-) diff --git a/KeyFinderLib/KeyFinder.cpp b/KeyFinderLib/KeyFinder.cpp index 19f56cbd..0ea45970 100644 --- a/KeyFinderLib/KeyFinder.cpp +++ b/KeyFinderLib/KeyFinder.cpp @@ -10,29 +10,27 @@ void KeyFinder::defaultResultCallback(KeySearchResult result) { - // Do nothing + // Do nothing } void KeyFinder::defaultStatusCallback(KeySearchStatus status) { - // Do nothing + // Do nothing } KeyFinder::KeyFinder(const secp256k1::uint256 &startKey, const secp256k1::uint256 &endKey, int compression, KeySearchDevice* device, const secp256k1::uint256 &stride) { - _total = 0; - _statusInterval = 1000; - _device = device; + _total = 0; + _statusInterval = 1000; + _device = device; - _compression = compression; + _compression = compression; _startKey = startKey; - _endKey = endKey; - _statusCallback = NULL; - - _resultCallback = NULL; + _statusCallback = NULL; + _resultCallback = NULL; _iterCount = 0; @@ -45,61 +43,61 @@ KeyFinder::~KeyFinder() void KeyFinder::setTargets(std::vector &targets) { - if(targets.size() == 0) { - throw KeySearchException("Requires at least 1 target"); - } + if(targets.size() == 0) { + throw KeySearchException("Requires at least 1 target"); + } - _targets.clear(); + _targets.clear(); - // Convert each address from base58 encoded form to a 160-bit integer - for(unsigned int i = 0; i < targets.size(); i++) { + // Convert each address from base58 encoded form to a 160-bit integer + for(unsigned int i = 0; i < targets.size(); i++) { - if(!Address::verifyAddress(targets[i])) { - throw KeySearchException("Invalid address '" + targets[i] + "'"); - } + if(!Address::verifyAddress(targets[i])) { + throw KeySearchException("Invalid address '" + targets[i] + "'"); + } - KeySearchTarget t; + KeySearchTarget t; - Base58::toHash160(targets[i], t.value); + Base58::toHash160(targets[i], t.value); - _targets.insert(t); - } + _targets.insert(t); + } _device->setTargets(_targets); } void KeyFinder::setTargets(std::string targetsFile) { - std::ifstream inFile(targetsFile.c_str()); + std::ifstream inFile(targetsFile.c_str()); - if(!inFile.is_open()) { - Logger::log(LogLevel::Error, "Unable to open '" + targetsFile + "'"); - throw KeySearchException(); - } + if(!inFile.is_open()) { + Logger::log(LogLevel::Error, "Unable to open '" + targetsFile + "'"); + throw KeySearchException(); + } - _targets.clear(); + _targets.clear(); - std::string line; - Logger::log(LogLevel::Info, "Loading addresses from '" + targetsFile + "'"); - while(std::getline(inFile, line)) { - util::removeNewline(line); + std::string line; + Logger::log(LogLevel::Info, "Loading addresses from '" + targetsFile + "'"); + while(std::getline(inFile, line)) { + util::removeNewline(line); line = util::trim(line); - if(line.length() > 0) { - if(!Address::verifyAddress(line)) { - Logger::log(LogLevel::Error, "Invalid address '" + line + "'"); - throw KeySearchException(); - } + if(line.length() > 0) { + if(!Address::verifyAddress(line)) { + Logger::log(LogLevel::Error, "Invalid address '" + line + "'"); + throw KeySearchException(); + } - KeySearchTarget t; + KeySearchTarget t; - Base58::toHash160(line, t.value); + Base58::toHash160(line, t.value); - _targets.insert(t); - } - } - Logger::log(LogLevel::Info, util::formatThousands(_targets.size()) + " addresses loaded (" - + util::format("%.1f", (double)(sizeof(KeySearchTarget) * _targets.size()) / (double)(1024 * 1024)) + "MB)"); + _targets.insert(t); + } + } + Logger::log(LogLevel::Info, util::formatThousands(_targets.size()) + " addresses loaded (" + + util::format("%.1f", (double)(sizeof(KeySearchTarget) * _targets.size()) / (double)(1024 * 1024)) + "MB)"); _device->setTargets(_targets); } @@ -107,34 +105,34 @@ void KeyFinder::setTargets(std::string targetsFile) void KeyFinder::setResultCallback(void(*callback)(KeySearchResult)) { - _resultCallback = callback; + _resultCallback = callback; } void KeyFinder::setStatusCallback(void(*callback)(KeySearchStatus)) { - _statusCallback = callback; + _statusCallback = callback; } void KeyFinder::setStatusInterval(uint64_t interval) { - _statusInterval = interval; + _statusInterval = interval; } void KeyFinder::setTargetsOnDevice() { - // Set the target in constant memory - std::vector targets; + // Set the target in constant memory + std::vector targets; - for(std::set::iterator i = _targets.begin(); i != _targets.end(); ++i) { - targets.push_back(hash160((*i).value)); - } + for(std::set::iterator i = _targets.begin(); i != _targets.end(); ++i) { + targets.push_back(hash160((*i).value)); + } _device->setTargets(_targets); } void KeyFinder::init() { - Logger::log(LogLevel::Info, "Initializing " + _device->getDeviceName()); + Logger::log(LogLevel::Info, "Initializing " + _device->getDeviceName()); _device->init(_startKey, _compression, _stride); } @@ -142,20 +140,20 @@ void KeyFinder::init() void KeyFinder::stop() { - _running = false; + _running = false; } void KeyFinder::removeTargetFromList(const unsigned int hash[5]) { - KeySearchTarget t(hash); + KeySearchTarget t(hash); - _targets.erase(t); + _targets.erase(t); } bool KeyFinder::isTargetInList(const unsigned int hash[5]) { - KeySearchTarget t(hash); - return _targets.find(t) != _targets.end(); + KeySearchTarget t(hash); + return _targets.find(t) != _targets.end(); } @@ -163,79 +161,79 @@ void KeyFinder::run() { uint64_t pointsPerIteration = _device->keysPerStep(); - _running = true; + _running = true; - util::Timer timer; + util::Timer timer; - timer.start(); + timer.start(); - uint64_t prevIterCount = 0; + uint64_t prevIterCount = 0; - _totalTime = 0; + _totalTime = 0; - while(_running) { + while(_running) { _device->doStep(); _iterCount++; - // Update status - uint64_t t = timer.getTime(); + // Update status + uint64_t t = timer.getTime(); - if(t >= _statusInterval) { + if(t >= _statusInterval) { - KeySearchStatus info; + KeySearchStatus info; - uint64_t count = (_iterCount - prevIterCount) * pointsPerIteration; + uint64_t count = (_iterCount - prevIterCount) * pointsPerIteration; - _total += count; + _total += count; - double seconds = (double)t / 1000.0; + double seconds = (double)t / 1000.0; - info.speed = (double)((double)count / seconds) / 1000000.0; + info.speed = (double)((double)count / seconds) / 1000000.0; - info.total = _total; + info.total = _total; - info.totalTime = _totalTime; + info.totalTime = _totalTime; - uint64_t freeMem = 0; + uint64_t freeMem = 0; - uint64_t totalMem = 0; + uint64_t totalMem = 0; - _device->getMemoryInfo(freeMem, totalMem); + _device->getMemoryInfo(freeMem, totalMem); - info.freeMemory = freeMem; - info.deviceMemory = totalMem; - info.deviceName = _device->getDeviceName(); - info.targets = _targets.size(); + info.freeMemory = freeMem; + info.deviceMemory = totalMem; + info.deviceName = _device->getDeviceName(); + info.targets = _targets.size(); info.nextKey = getNextKey(); - _statusCallback(info); + _statusCallback(info); - timer.start(); - prevIterCount = _iterCount; - _totalTime += t; - } + timer.start(); + prevIterCount = _iterCount; + _totalTime += t; + } std::vector results; if(_device->getResults(results) > 0) { - for(unsigned int i = 0; i < results.size(); i++) { + for(unsigned int i = 0; i < results.size(); i++) { - KeySearchResult info; + KeySearchResult info; info.privateKey = results[i].privateKey; info.publicKey = results[i].publicKey; - info.compressed = results[i].compressed; - info.address = Address::fromPublicKey(results[i].publicKey, results[i].compressed); + info.compressed = results[i].compressed; + info.address = Address::fromPublicKey(results[i].publicKey, results[i].compressed); - _resultCallback(info); - } + _resultCallback(info); + } - // Remove the hashes that were found - for(unsigned int i = 0; i < results.size(); i++) { - removeTargetFromList(results[i].hash); - } - } + // Remove the hashes that were found + for(unsigned int i = 0; i < results.size(); i++) { + removeTargetFromList(results[i].hash); + } + } // Stop if there are no keys left if(_targets.size() == 0) { @@ -243,12 +241,12 @@ void KeyFinder::run() _running = false; } - // Stop if we searched the entire range + // Stop if we searched the entire range if(_device->getNextKey().cmp(_endKey) >= 0 || _device->getNextKey().cmp(_startKey) < 0) { Logger::log(LogLevel::Info, "Reached end of keyspace"); _running = false; } - } + } } secp256k1::uint256 KeyFinder::getNextKey() diff --git a/KeyFinderLib/KeyFinder.h b/KeyFinderLib/KeyFinder.h index c8150f7d..dd7e5ae7 100644 --- a/KeyFinderLib/KeyFinder.h +++ b/KeyFinderLib/KeyFinder.h @@ -15,50 +15,50 @@ class KeyFinder { KeySearchDevice *_device; - unsigned int _compression; + unsigned int _compression; - std::set _targets; + std::set _targets; - uint64_t _statusInterval; + uint64_t _statusInterval; secp256k1::uint256 _stride = 1; - uint64_t _iterCount; - uint64_t _total; - uint64_t _totalTime; + uint64_t _iterCount; + uint64_t _total; + uint64_t _totalTime; secp256k1::uint256 _startKey; secp256k1::uint256 _endKey; - // Each index of each thread gets a flag to indicate if it found a valid hash - bool _running; + // Each index of each thread gets a flag to indicate if it found a valid hash + bool _running; - void(*_resultCallback)(KeySearchResult); - void(*_statusCallback)(KeySearchStatus); + void(*_resultCallback)(KeySearchResult); + void(*_statusCallback)(KeySearchStatus); - static void defaultResultCallback(KeySearchResult result); - static void defaultStatusCallback(KeySearchStatus status); + static void defaultResultCallback(KeySearchResult result); + static void defaultStatusCallback(KeySearchStatus status); - void removeTargetFromList(const unsigned int value[5]); - bool isTargetInList(const unsigned int value[5]); - void setTargetsOnDevice(); + void removeTargetFromList(const unsigned int value[5]); + bool isTargetInList(const unsigned int value[5]); + void setTargetsOnDevice(); public: KeyFinder(const secp256k1::uint256 &startKey, const secp256k1::uint256 &endKey, int compression, KeySearchDevice* device, const secp256k1::uint256 &stride); - ~KeyFinder(); + ~KeyFinder(); - void init(); - void run(); - void stop(); + void init(); + void run(); + void stop(); - void setResultCallback(void(*callback)(KeySearchResult)); - void setStatusCallback(void(*callback)(KeySearchStatus)); - void setStatusInterval(uint64_t interval); + void setResultCallback(void(*callback)(KeySearchResult)); + void setStatusCallback(void(*callback)(KeySearchStatus)); + void setStatusInterval(uint64_t interval); - void setTargets(std::string targetFile); - void setTargets(std::vector &targets); + void setTargets(std::string targetFile); + void setTargets(std::vector &targets); secp256k1::uint256 getNextKey(); }; diff --git a/KeyFinderLib/KeyFinderShared.h b/KeyFinderLib/KeyFinderShared.h index a6973f1a..bd111486 100644 --- a/KeyFinderLib/KeyFinderShared.h +++ b/KeyFinderLib/KeyFinderShared.h @@ -2,22 +2,22 @@ #define _KEY_FINDER_SHARED_H namespace PointCompressionType { - enum Value { - COMPRESSED = 0, - UNCOMPRESSED = 1, - BOTH = 2 - }; + enum Value { + COMPRESSED = 0, + UNCOMPRESSED = 1, + BOTH = 2 + }; } // Structures that exist on both host and device side struct KeyFinderDeviceResult { - int thread; - int block; - int idx; - bool compressed; - unsigned int x[8]; - unsigned int y[8]; - unsigned int digest[5]; + int thread; + int block; + int idx; + bool compressed; + unsigned int x[8]; + unsigned int y[8]; + unsigned int digest[5]; }; //typedef struct hash160 { diff --git a/util/util.cpp b/util/util.cpp index b48750b7..4f350772 100644 --- a/util/util.cpp +++ b/util/util.cpp @@ -53,227 +53,227 @@ namespace util { #endif } - std::string formatThousands(uint64_t x) - { - char buf[32] = ""; - - sprintf(buf, "%lld", x); - - std::string s(buf); - - int len = (int)s.length(); + std::string formatThousands(uint64_t x) + { + char buf[32] = ""; - int numCommas = (len - 1) / 3; + sprintf(buf, "%lld", x); - if(numCommas == 0) { - return s; - } + std::string s(buf); - std::string result = ""; + int len = (int)s.length(); - int count = ((len % 3) == 0) ? 0 : (3 - (len % 3)); + int numCommas = (len - 1) / 3; - for(int i = 0; i < len; i++) { - result += s[i]; + if(numCommas == 0) { + return s; + } - if(count++ == 2 && i < len - 1) { - result += ","; - count = 0; - } - } + std::string result = ""; - return result; - } + int count = ((len % 3) == 0) ? 0 : (3 - (len % 3)); - uint32_t parseUInt32(std::string s) - { - return (uint32_t)parseUInt64(s); - } + for(int i = 0; i < len; i++) { + result += s[i]; - uint64_t parseUInt64(std::string s) - { - uint64_t val = 0; - bool isHex = false; + if(count++ == 2 && i < len - 1) { + result += ","; + count = 0; + } + } - if(s[0] == '0' && s[1] == 'x') { - isHex = true; - s = s.substr(2); - } - - if(s[s.length() - 1] == 'h') { - isHex = true; - s = s.substr(0, s.length() - 1); - } + return result; + } - if(isHex) { - if(sscanf(s.c_str(), "%llx", &val) != 1) { - throw std::string("Expected an integer"); - } - } else { - if(sscanf(s.c_str(), "%lld", &val) != 1) { - throw std::string("Expected an integer"); - } - } + uint32_t parseUInt32(std::string s) + { + return (uint32_t)parseUInt64(s); + } - return val; - } + uint64_t parseUInt64(std::string s) + { + uint64_t val = 0; + bool isHex = false; + + if(s[0] == '0' && s[1] == 'x') { + isHex = true; + s = s.substr(2); + } + + if(s[s.length() - 1] == 'h') { + isHex = true; + s = s.substr(0, s.length() - 1); + } + + if(isHex) { + if(sscanf(s.c_str(), "%llx", &val) != 1) { + throw std::string("Expected an integer"); + } + } else { + if(sscanf(s.c_str(), "%lld", &val) != 1) { + throw std::string("Expected an integer"); + } + } + + return val; + } - bool isHex(const std::string &s) - { - int len = 0; + bool isHex(const std::string &s) + { + int len = 0; - for(int i = 0; i < len; i++) { - char c = s[i]; + for(int i = 0; i < len; i++) { + char c = s[i]; - if(!((c >= '0' && c <= '9') || (c >= 'a' && c <= 'f') || (c >= 'A' && c <= 'F'))) { - return false; - } - } + if(!((c >= '0' && c <= '9') || (c >= 'a' && c <= 'f') || (c >= 'A' && c <= 'F'))) { + return false; + } + } - return true; - } + return true; + } - std::string formatSeconds(unsigned int seconds) - { - char s[128] = { 0 }; + std::string formatSeconds(unsigned int seconds) + { + char s[128] = { 0 }; - unsigned int days = seconds / 86400; - unsigned int hours = (seconds % 86400) / 3600; - unsigned int minutes = (seconds % 3600) / 60; - unsigned int sec = seconds % 60; + unsigned int days = seconds / 86400; + unsigned int hours = (seconds % 86400) / 3600; + unsigned int minutes = (seconds % 3600) / 60; + unsigned int sec = seconds % 60; - if(days > 0) { - sprintf(s, "%d:%02d:%02d:%02d", days, hours, minutes, sec); - } else { - sprintf(s, "%02d:%02d:%02d", hours, minutes, sec); - } - + if(days > 0) { + sprintf(s, "%d:%02d:%02d:%02d", days, hours, minutes, sec); + } else { + sprintf(s, "%02d:%02d:%02d", hours, minutes, sec); + } + - return std::string(s); - } + return std::string(s); + } - long getFileSize(const std::string &fileName) - { - FILE *fp = fopen(fileName.c_str(), "rb"); - if(fp == NULL) { - return -1; - } + long getFileSize(const std::string &fileName) + { + FILE *fp = fopen(fileName.c_str(), "rb"); + if(fp == NULL) { + return -1; + } - fseek(fp, 0, SEEK_END); + fseek(fp, 0, SEEK_END); - long pos = ftell(fp); + long pos = ftell(fp); - fclose(fp); + fclose(fp); - return pos; - } + return pos; + } - bool readLinesFromStream(const std::string &fileName, std::vector &lines) - { - std::ifstream inFile(fileName.c_str()); - - if(!inFile.is_open()) { - return false; - } - - return readLinesFromStream(inFile, lines); - } - - bool readLinesFromStream(std::istream &in, std::vector &lines) - { - std::string line; - - while(std::getline(in, line)) { - if(line.length() > 0) { - lines.push_back(line); - } - } - - return true; - } - - bool appendToFile(const std::string &fileName, const std::string &s) - { - std::ofstream outFile; - bool newline = false; - - if(getFileSize(fileName) > 0) { - newline = true; - } + bool readLinesFromStream(const std::string &fileName, std::vector &lines) + { + std::ifstream inFile(fileName.c_str()); - outFile.open(fileName.c_str(), std::ios::app); + if(!inFile.is_open()) { + return false; + } - if(!outFile.is_open()) { - return false; - } + return readLinesFromStream(inFile, lines); + } - // Add newline following previous line - if(newline) { - outFile << std::endl; - } + bool readLinesFromStream(std::istream &in, std::vector &lines) + { + std::string line; - outFile << s; + while(std::getline(in, line)) { + if(line.length() > 0) { + lines.push_back(line); + } + } - return true; - } + return true; + } - std::string format(const char *formatStr, double value) - { - char buf[100] = { 0 }; + bool appendToFile(const std::string &fileName, const std::string &s) + { + std::ofstream outFile; + bool newline = false; - sprintf(buf, formatStr, value); + if(getFileSize(fileName) > 0) { + newline = true; + } - return std::string(buf); - } + outFile.open(fileName.c_str(), std::ios::app); - std::string format(uint32_t value) - { - char buf[100] = { 0 }; + if(!outFile.is_open()) { + return false; + } - sprintf(buf, "%u", value); + // Add newline following previous line + if(newline) { + outFile << std::endl; + } - return std::string(buf); - } + outFile << s; - std::string format(uint64_t value) - { - char buf[100] = { 0 }; + return true; + } + + std::string format(const char *formatStr, double value) + { + char buf[100] = { 0 }; + + sprintf(buf, formatStr, value); + + return std::string(buf); + } + + std::string format(uint32_t value) + { + char buf[100] = { 0 }; - sprintf(buf, "%lld", (uint64_t)value); + sprintf(buf, "%u", value); - return std::string(buf); - } + return std::string(buf); + } - std::string format(int value) - { - char buf[100] = { 0 }; + std::string format(uint64_t value) + { + char buf[100] = { 0 }; - sprintf(buf, "%d", value); + sprintf(buf, "%lld", (uint64_t)value); - return std::string(buf); - } + return std::string(buf); + } - void removeNewline(std::string &s) - { - size_t len = s.length(); + std::string format(int value) + { + char buf[100] = { 0 }; - int toRemove = 0; + sprintf(buf, "%d", value); - if(len >= 2) { - if(s[len - 2] == '\r' || s[len - 2] == '\n') { - toRemove++; - } - } - if(len >= 1) { - if(s[len - 1] == '\r' || s[len - 1] == '\n') { - toRemove++; - } - } + return std::string(buf); + } - if(toRemove) { - s.erase(len - toRemove); - } - } + void removeNewline(std::string &s) + { + size_t len = s.length(); + + int toRemove = 0; + + if(len >= 2) { + if(s[len - 2] == '\r' || s[len - 2] == '\n') { + toRemove++; + } + } + if(len >= 1) { + if(s[len - 1] == '\r' || s[len - 1] == '\n') { + toRemove++; + } + } + + if(toRemove) { + s.erase(len - toRemove); + } + } unsigned int endian(unsigned int x) {