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/CudaDeviceKeys.cu b/CudaKeySearchDevice/CudaDeviceKeys.cu index d98dbaa7..e4183ad6 100644 --- a/CudaKeySearchDevice/CudaDeviceKeys.cu +++ b/CudaKeySearchDevice/CudaDeviceKeys.cu @@ -22,8 +22,22 @@ __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); +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) { @@ -45,15 +59,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 += totalThreads; + index++; + } + + index = base + totalThreads * 4 + threadId; + + for(int k = 4; k < 8; k++) { + dest[index] = value[k]; + index++; } } @@ -62,15 +83,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++; + } + + index = base + totalThreads * 4 + threadId; + + for(int k = 4; k < 8; k++) { value[k] = src[index]; - index += totalThreads; + index++; } secp256k1::uint256 v(value, secp256k1::uint256::BigEndian); @@ -255,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; } @@ -267,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; @@ -286,7 +316,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 +424,4 @@ bool CudaDeviceKeys::selfTest(const std::vector &privateKeys } return true; -} \ No newline at end of file +} 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/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/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/KeyFinder/main.cpp b/KeyFinder/main.cpp index 22c699df..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) @@ -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; - } - - - 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); + 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); 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(); 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/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 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. 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 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 diff --git a/cudaMath/sha256.cuh b/cudaMath/sha256.cuh index b04aa9a0..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 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 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; @@ -98,22 +96,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 +132,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 +167,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 +204,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 +244,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 +280,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 +314,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 +348,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 +402,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 +436,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 +470,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 +505,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 +540,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 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) {