Skip to content

Commit

Permalink
Bug fix + code cleanup
Browse files Browse the repository at this point in the history 
-Moved CudaDeviceContext into its own header/source seperate from DeviceContext
-Store target hashes in std::set instead of std::vector
-Fixed bug where BitCracker.exe throw an exception when both compressed and uncompressed is enabled, and multiple results are found in the same kernel call and some are compressed and some are uncompressed
  • Loading branch information
@brichard19
brichard19 committed Jul 28, 2018
1 parent d8774f6 commit f06540d
Show file tree
Hide file tree
Showing 10 changed files with 220 additions and 174 deletions.
24 changes: 0 additions & 24 deletions CmdParse/CmdParse.cpp
View file
Original file line number Diff line number Diff line change
@@ -1,29 +1,5 @@
#include "CmdParse.h"


//static std::vector<std::string> splitString(const std::string s, const std::string delimeter)
//{
// std::vector<int> idxVec;
//
// int idx = 0;
// while((idx = s.find_first_of(delimeter, idx)) != -1) {
// idxVec.push_back(idx);
// }
//
// std::vector<std::string> 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()
{

Expand Down
100 changes: 48 additions & 52 deletions KeyFinderLib/KeyFinder.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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)
{
Expand All @@ -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) {
Expand All @@ -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;
}

Expand Down Expand Up @@ -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<struct hash160> 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<KeyFinderTarget>::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);
}
}

Expand All @@ -139,14 +142,15 @@ void KeyFinder::init()
generateStartingPoints();
_devCtx->copyPoints(_startingPoints);

setTargetHashes();
setTargetsOnDevice();

allocateChainBuf(_numThreads * _numBlocks * _pointsPerThread);

// Set the incrementor
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));
Expand All @@ -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");
}
Expand Down Expand Up @@ -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<KeyFinderTarget>::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<KeyFinderTarget>::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<KeyFinderResult> &r)
Expand All @@ -260,7 +255,7 @@ void KeyFinder::getResults(std::vector<KeyFinderResult> &r)
struct KeyFinderDeviceResult *rPtr = &((struct KeyFinderDeviceResult *)ptr)[i];

// might be false-positive
if(!isHashInList(rPtr->digest)) {
if(!isTargetInList(rPtr->digest)) {
continue;
}

Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -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");
}

Expand All @@ -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;
}
Expand Down
46 changes: 33 additions & 13 deletions KeyFinderLib/KeyFinder.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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)
{
Expand All @@ -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
Expand Down Expand Up @@ -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);
}
Expand All @@ -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));
}

Expand All @@ -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;
}
}
Expand Down Expand Up @@ -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);
Expand All @@ -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)) {
Expand Down
Loading

0 comments on commit f06540d

Please sign in to comment.