diff --git a/CmdParse/CmdParse.cpp b/CmdParse/CmdParse.cpp index 1418a74..3cd6c49 100644 --- a/CmdParse/CmdParse.cpp +++ b/CmdParse/CmdParse.cpp @@ -1,29 +1,5 @@ #include "CmdParse.h" - -//static std::vector splitString(const std::string s, const std::string delimeter) -//{ -// std::vector idxVec; -// -// int idx = 0; -// while((idx = s.find_first_of(delimeter, idx)) != -1) { -// idxVec.push_back(idx); -// } -// -// std::vector vec; -// -// int prev = 0; -// for(int i = 0; i < idxVec.size(); i++) { -// -// std::string tmp = s.substr(prev, idxVec[i] - prev); -// prev = idxVec[i]; -// -// vec.push_back(tmp); -// } -// -// return vec; -//} - CmdParse::CmdParse() { diff --git a/KeyFinderLib/KeyFinder.cpp b/KeyFinderLib/KeyFinder.cpp index b540c8d..db5982d 100644 --- a/KeyFinderLib/KeyFinder.cpp +++ b/KeyFinderLib/KeyFinder.cpp @@ -2,13 +2,9 @@ #include "util.h" #include "AddressUtil.h" -#include "DeviceContext.h" +#include "CudaDeviceContext.h" #include "cudabridge.h" -static const int DEFAULT_POINTS_PER_THREAD = 1; -static const int DEFAULT_NUM_THREADS = 32; -static const int DEFAULT_NUM_BLOCKS = 1; - void KeyFinder::defaultResultCallback(KeyFinderResultInfo result) { @@ -27,25 +23,22 @@ KeyFinder::KeyFinder(int device, const secp256k1::uint256 &start, unsigned long _total = 0; _statusInterval = 1000; _device = device; - _pointsPerThread = DEFAULT_POINTS_PER_THREAD; - _numThreads = DEFAULT_NUM_THREADS; - _numBlocks = DEFAULT_NUM_BLOCKS; - if(!(compression == Compression::COMPRESSED || compression == Compression::UNCOMPRESSED || compression == Compression::BOTH)) { - throw KeyFinderException("Invalid argument for compression"); + + if(threads <= 0 || threads % 32 != 0) { + throw KeyFinderException("The number of threads must be a multiple of 32"); } - _compression = compression; - if(threads != 0) { - _numThreads = threads; + if(blocks <= 0) { + throw KeyFinderException("At least one block required"); } - if(blocks != 0) { - _numBlocks = blocks; + if(pointsPerThread <= 0) { + throw KeyFinderException("At least one point per thread required"); } - if(pointsPerThread != 0) { - _pointsPerThread = pointsPerThread; + if(!(compression == Compression::COMPRESSED || compression == Compression::UNCOMPRESSED || compression == Compression::BOTH)) { + throw KeyFinderException("Invalid argument for compression"); } if(start.cmp(secp256k1::N) >= 0) { @@ -58,21 +51,32 @@ KeyFinder::KeyFinder(int device, const secp256k1::uint256 &start, unsigned long // Convert each address from base58 encoded form to a 160-bit integer for(unsigned int i = 0; i < targetHashes.size(); i++) { - KeyFinderTarget t; - + if(!Address::verifyAddress(targetHashes[i])) { - throw KeyFinderException("Invalid address"); + throw KeyFinderException("Invalid address '" + targetHashes[i] + "'"); } - Base58::toHash160(targetHashes[i], t.hash); + KeyFinderTarget t; + + Base58::toHash160(targetHashes[i], t.value); - _targets.push_back(t); + _targets.insert(t); } + _compression = compression; + + _numThreads = threads; + + _numBlocks = blocks; + + _pointsPerThread = pointsPerThread; + _startExponent = start; + _range = range; _statusCallback = NULL; + _resultCallback = NULL; } @@ -102,21 +106,20 @@ void KeyFinder::setStatusInterval(unsigned int interval) _statusInterval = interval; } -void KeyFinder::setTargetHashes() +void KeyFinder::setTargetsOnDevice() { // Set the target in constant memory std::vector targets; - for(int i = 0; i < _targets.size(); i++) { - struct hash160 h; - memcpy(h.h, _targets[i].hash, sizeof(unsigned int) * 5); - targets.push_back(h); + + for(std::set::iterator i = _targets.begin(); i != _targets.end(); ++i) { + targets.push_back(hash160((*i).value)); } cudaError_t err = setTargetHash(targets); if(err) { std::string cudaErrorString(cudaGetErrorString(err)); - throw KeyFinderException("Error initializing device: " + cudaErrorString); + throw KeyFinderException("Device error: " + cudaErrorString); } } @@ -139,7 +142,7 @@ void KeyFinder::init() generateStartingPoints(); _devCtx->copyPoints(_startingPoints); - setTargetHashes(); + setTargetsOnDevice(); allocateChainBuf(_numThreads * _numBlocks * _pointsPerThread); @@ -147,6 +150,7 @@ void KeyFinder::init() secp256k1::ecpoint g = secp256k1::G(); secp256k1::ecpoint p = secp256k1::multiplyPoint(secp256k1::uint256(_numThreads * _numBlocks * _pointsPerThread), g); + cudaError_t err = setIncrementorPoint(p.x, p.y); if(err) { std::string cudaErrorString(cudaGetErrorString(err)); @@ -173,7 +177,7 @@ void KeyFinder::generateStartingPoints() secp256k1::generateKeyPairsBulk(secp256k1::G(), _exponents, _startingPoints); - for(unsigned long long i = 0; i < _startingPoints.size(); i++) { + for(unsigned int i = 0; i < _startingPoints.size(); i++) { if(!secp256k1::pointExists(_startingPoints[i])) { throw KeyFinderException("Point generation error"); } @@ -222,25 +226,16 @@ bool KeyFinder::verifyKey(const secp256k1::uint256 &privateKey, const secp256k1: return true; } -void KeyFinder::removeHashFromList(const unsigned int hash[5]) +void KeyFinder::removeTargetFromList(const unsigned int hash[5]) { - for(std::vector::iterator i = _targets.begin(); i != _targets.end(); ++i) { - if(memcmp((*i).hash, hash, sizeof(unsigned int) * 5) == 0) { - _targets.erase(i); - break; - } - } + KeyFinderTarget t(hash); + _targets.erase(t); } -bool KeyFinder::isHashInList(const unsigned int hash[5]) +bool KeyFinder::isTargetInList(const unsigned int hash[5]) { - for(std::vector::iterator i = _targets.begin(); i != _targets.end(); ++i) { - if(memcmp((*i).hash, hash, sizeof(unsigned int) * 5) == 0) { - return true; - } - } - - return false; + KeyFinderTarget t(hash); + return _targets.find(t) != _targets.end(); } void KeyFinder::getResults(std::vector &r) @@ -260,7 +255,7 @@ void KeyFinder::getResults(std::vector &r) struct KeyFinderDeviceResult *rPtr = &((struct KeyFinderDeviceResult *)ptr)[i]; // might be false-positive - if(!isHashInList(rPtr->digest)) { + if(!isTargetInList(rPtr->digest)) { continue; } @@ -294,6 +289,8 @@ void KeyFinder::run() while(_running) { + _devCtx->clearResults(); + KernelParams params = _devCtx->getKernelParams(); if(_iterCount < 2 && _startExponent.cmp(pointsPerIteration) <= 0) { callKeyFinderKernel(params, true, _compression); @@ -337,7 +334,7 @@ void KeyFinder::run() unsigned long long offset = (unsigned long long)_numBlocks * _numThreads * _pointsPerThread * _iterCount; exp = secp256k1::addModN(exp, secp256k1::uint256(offset)); - if(!verifyKey(exp, publicKey, results[i].hash, results[0].compressed)) { + if(!verifyKey(exp, publicKey, results[i].hash, results[i].compressed)) { throw KeyFinderException("Invalid point"); } @@ -352,17 +349,16 @@ void KeyFinder::run() // Remove the hashes that were found - for(int i = 0; i < results.size(); i++) { - removeHashFromList(results[i].hash); + for(unsigned int i = 0; i < results.size(); i++) { + removeTargetFromList(results[i].hash); } // Update hash targets on device - setTargetHashes(); - - //_running = false; + setTargetsOnDevice(); } _iterCount++; + // Stop if we searched the entire range, or have no targets left if((_range > 0 && _iterCount * pointsPerIteration >= _range) || _targets.size() == 0) { _running = false; } diff --git a/KeyFinderLib/KeyFinder.cu b/KeyFinderLib/KeyFinder.cu index 354b8ff..60b6471 100644 --- a/KeyFinderLib/KeyFinder.cu +++ b/KeyFinderLib/KeyFinder.cu @@ -29,13 +29,7 @@ static bool _useBloomFilter = false; static unsigned int *_bloomFilterPtr = NULL; static unsigned int *_chainBufferPtr = NULL; -static const unsigned int _RIPEMD160_IV_HOST[5] = { - 0x67452301, - 0xefcdab89, - 0x98badcfe, - 0x10325476, - 0xc3d2e1f0 -}; + static unsigned int swp(unsigned int x) { @@ -45,11 +39,33 @@ static unsigned int swp(unsigned int x) static void undoRMD160FinalRound(const unsigned int hIn[5], unsigned int hOut[5]) { + unsigned int iv[5] = { + 0x67452301, + 0xefcdab89, + 0x98badcfe, + 0x10325476, + 0xc3d2e1f0 + }; + for(int i = 0; i < 5; i++) { - hOut[i] = swp(hIn[i]) - _RIPEMD160_IV_HOST[(i + 1) % 5]; + hOut[i] = swp(hIn[i]) - iv[(i + 1) % 5]; } } +__device__ void doRMD160FinalRound(const unsigned int hIn[5], unsigned int hOut[5]) +{ + 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]); + } +} /** Copies the target hashes to constant memory @@ -245,6 +261,7 @@ __device__ void hashPublicKeyCompressed(const unsigned int *x, unsigned int yPar __device__ void addResult(unsigned int *numResultsPtr, void *results, void *info, unsigned int size) { unsigned int count = atomicAdd(numResultsPtr, 1); + unsigned char *ptr = (unsigned char *)results + count * size; memcpy(ptr, info, size); } @@ -263,9 +280,11 @@ __device__ void setResultFound(unsigned int *numResultsPtr, void *results, int i r.y[i] = y[i]; } - for(int i = 0; i < 5; i++) { - r.digest[i] = endian(digest[i] + _RIPEMD160_IV[(i + 1) % 5]); - } + //for(int i = 0; i < 5; i++) { + // r.digest[i] = endian(digest[i] + _RIPEMD160_IV[(i + 1) % 5]); + //} + doRMD160FinalRound(digest, r.digest); + addResult(numResultsPtr, results, &r, sizeof(r)); } @@ -289,6 +308,7 @@ __device__ bool checkHash(unsigned int hash[5]) for(int i = 0; i < 5; i++) { equal &= (hash[i] == _TARGET_HASH[j][i]); } + foundMatch |= equal; } } @@ -362,7 +382,7 @@ __device__ void doIterationWithDouble(unsigned int *xPtr, unsigned int *yPtr, in // uncompressed - if(compression == 1 || compression == 2) { + if(compression == PointCompressionType::UNCOMPRESSED || compression == PointCompressionType::BOTH) { unsigned int y[8]; readInt(yPtr, i, y); hashPublicKey(x, y, digest); @@ -373,7 +393,7 @@ __device__ void doIterationWithDouble(unsigned int *xPtr, unsigned int *yPtr, in } // compressed - if(compression == 0 || compression == 2) { + if(compression == PointCompressionType::COMPRESSED || compression == PointCompressionType::BOTH) { hashPublicKeyCompressed(x, readIntLSW(yPtr, i), digest); if(checkHash(digest)) { diff --git a/KeyFinderLib/KeyFinder.h b/KeyFinderLib/KeyFinder.h index 4067f89..7d3aac9 100644 --- a/KeyFinderLib/KeyFinder.h +++ b/KeyFinderLib/KeyFinder.h @@ -2,11 +2,22 @@ #define _KEY_FINDER_H #include +#include #include "secp256k1.h" class CudaDeviceContext; +struct KeyFinderResult { + int thread; + int block; + int index; + bool compressed; + + secp256k1::ecpoint p; + unsigned int hash[5]; +}; + typedef struct { std::string address; secp256k1::ecpoint publicKey; @@ -25,9 +36,59 @@ typedef struct { }KeyFinderStatusInfo; -struct KeyFinderTarget { - secp256k1::ecpoint p; - unsigned int hash[5]; +class KeyFinderTarget { + +public: + unsigned int value[5] = { 0 }; + + KeyFinderTarget() + { + } + + KeyFinderTarget(const unsigned int h[5]) + { + for(int i = 0; i < 5; i++) { + value[i] = h[i]; + } + } + + + bool operator==(const KeyFinderTarget &t) const + { + for(int i = 0; i < 5; i++) { + if(value[i] != t.value[i]) { + return false; + } + } + + return true; + } + + bool operator<(const KeyFinderTarget &t) const + { + for(int i = 0; i < 5; i++) { + if(value[i] < t.value[i]) { + return true; + } else if(value[i] > t.value[i]) { + return false; + } + } + + return false; + } + + bool operator>(const KeyFinderTarget &t) const + { + for(int i = 0; i < 5; i++) { + if(value[i] > t.value[i]) { + return true; + } else if(value[i] < t.value[i]) { + return false; + } + } + + return false; + } }; class KeyFinderException { @@ -42,19 +103,15 @@ class KeyFinderException { std::string msg; }; -class KeyFinderResult; - class KeyFinder { private: - static const int FORMAT_PUBKEY = 0x00000001; - unsigned int _compression; unsigned int _flags; - std::vector _targets; + std::set _targets; unsigned int _statusInterval; @@ -98,9 +155,9 @@ class KeyFinder { void getResults(std::vector &r); - void removeHashFromList(const unsigned int hash[5]); - bool isHashInList(const unsigned int hash[5]); - void setTargetHashes(); + void removeTargetFromList(const unsigned int value[5]); + bool isTargetInList(const unsigned int value[5]); + void setTargetsOnDevice(); public: class Compression { diff --git a/cudaDeviceContext/DeviceContext.cpp b/cudaDeviceContext/CudaDeviceContext.cpp similarity index 95% rename from cudaDeviceContext/DeviceContext.cpp rename to cudaDeviceContext/CudaDeviceContext.cpp index 5995e31..0fe5aef 100644 --- a/cudaDeviceContext/DeviceContext.cpp +++ b/cudaDeviceContext/CudaDeviceContext.cpp @@ -1,6 +1,9 @@ #include -#include "DeviceContext.h" +#include +#include + +#include "CudaDeviceContext.h" static std::string getErrorString(cudaError_t err) @@ -158,21 +161,13 @@ void CudaDeviceContext::splatBigInt(unsigned int *dest, int block, int thread, i void CudaDeviceContext::cleanup() { - if(_x != NULL) { - cudaFree(_y); - } + cudaFree(_y); - if(_y != NULL) { - cudaFree(_y); - } + cudaFree(_y); - if(_numResultsHost != NULL) { - cudaFreeHost(_numResultsHost); - } + cudaFreeHost(_numResultsHost); - if(_resultsHost != NULL) { - cudaFreeHost(_resultsHost); - } + cudaFreeHost(_resultsHost); _x = NULL; _y = NULL; diff --git a/cudaDeviceContext/CudaDeviceContext.h b/cudaDeviceContext/CudaDeviceContext.h new file mode 100644 index 0000000..cd16b3c --- /dev/null +++ b/cudaDeviceContext/CudaDeviceContext.h @@ -0,0 +1,53 @@ +#ifndef _CUDA_DEVICE_CONTEXT_H +#define _CUDA_DEVICE_CONTEXT_H + +#include "DeviceContext.h" + +class CudaDeviceContext : DeviceContext { + +private: + int _device; + + int _threads; + int _blocks; + int _pointsPerThread; + + unsigned int *_x; + unsigned int *_y; + + unsigned int *_numResultsHost; + unsigned int *_numResultsDev; + unsigned int *_resultsHost; + unsigned int *_resultsDev; + + void splatBigInt(unsigned int *dest, int block, int thread, int idx, const secp256k1::uint256 &value); + +public: + CudaDeviceContext() + { + _device = 0; + _threads = 0; + _blocks = 0; + _pointsPerThread = 0; + + _x = NULL; + _y = NULL; + } + + void init(const DeviceParameters ¶ms); + + void copyPoints(const std::vector &points); + int getIndex(int block, int thread, int idx); + + KernelParams getKernelParams(); + + void cleanup(); + + bool resultFound(); + int getResultCount(); + void getResults(void *ptr, int size); + void clearResults(); + ~CudaDeviceContext(); +}; + +#endif \ No newline at end of file diff --git a/cudaDeviceContext/DeviceContext.h b/cudaDeviceContext/DeviceContext.h index 7696859..e9834eb 100644 --- a/cudaDeviceContext/DeviceContext.h +++ b/cudaDeviceContext/DeviceContext.h @@ -1,22 +1,11 @@ #ifndef _DEVICE_CONTEXT_H #define _DEVICE_CONTEXT_H -#include -#include + #include #include "secp256k1.h" #include "DeviceContextShared.h" -struct KeyFinderResult { - int thread; - int block; - int index; - bool compressed; - - secp256k1::ecpoint p; - unsigned int hash[5]; -}; - typedef struct { int blocks; int threads; @@ -68,51 +57,4 @@ class DeviceContext { virtual ~DeviceContext() {} }; -class CudaDeviceContext : DeviceContext { - -private: - int _device; - - int _threads; - int _blocks; - int _pointsPerThread; - - unsigned int *_x; - unsigned int *_y; - - unsigned int *_numResultsHost; - unsigned int *_numResultsDev; - unsigned int *_resultsHost; - unsigned int *_resultsDev; - - void splatBigInt(unsigned int *dest, int block, int thread, int idx, const secp256k1::uint256 &value); - -public: - CudaDeviceContext() - { - _device = 0; - _threads = 0; - _blocks = 0; - _pointsPerThread = 0; - - _x = NULL; - _y = NULL; - } - - void init(const DeviceParameters ¶ms); - - void copyPoints(const std::vector &points); - int getIndex(int block, int thread, int idx); - - KernelParams getKernelParams(); - - void cleanup(); - - bool resultFound(); - int getResultCount(); - void getResults(void *ptr, int size); - void clearResults(); - ~CudaDeviceContext(); -}; - #endif \ No newline at end of file diff --git a/cudaDeviceContext/DeviceContextShared.h b/cudaDeviceContext/DeviceContextShared.h index f46347e..f65f426 100644 --- a/cudaDeviceContext/DeviceContextShared.h +++ b/cudaDeviceContext/DeviceContextShared.h @@ -13,7 +13,13 @@ struct KeyFinderDeviceResult { }; struct hash160 { + unsigned int h[5]; + + hash160(const unsigned int hash[5]) + { + memcpy(h, hash, sizeof(unsigned int) * 5); + } }; namespace PointCompressionType { diff --git a/cudaDeviceContext/cudaDeviceContext.vcxproj b/cudaDeviceContext/cudaDeviceContext.vcxproj index 5dec6ca..7b80b89 100644 --- a/cudaDeviceContext/cudaDeviceContext.vcxproj +++ b/cudaDeviceContext/cudaDeviceContext.vcxproj @@ -19,9 +19,10 @@ - + + diff --git a/cudaMath/secp256k1.cuh b/cudaMath/secp256k1.cuh index b433857..8c4d5d4 100644 --- a/cudaMath/secp256k1.cuh +++ b/cudaMath/secp256k1.cuh @@ -55,7 +55,7 @@ __device__ __forceinline__ void copyBigInt(const unsigned int src[8], unsigned i } -__device__ void printBigInt(const unsigned int *x, int len) +__host__ __device__ void printBigInt(const unsigned int *x, int len) { for(int i = 0; i < len; i++) { printf("%.8x", x[i]);