Hex parsing bug fix

-Fixed hex string parsing, was not accepting 0x prefix or 'h' suffix -Removed unused kernel parameter 'flags' -When reading results from the device, DeviceContext::getResults() takes a void pointer instead of a struct, to make it more general.
Slyfrosty · Jun 14, 2018 · 88ab855 · 88ab855
1 parent c7a097e
commit 88ab855
Show file tree

Hide file tree

Showing 10 changed files with 155 additions and 88 deletions.
diff --git a/KeyFinderLib/KeyFinder.cpp b/KeyFinderLib/KeyFinder.cpp
@@ -103,8 +103,15 @@ void KeyFinder::setStatusInterval(unsigned int interval)
 void KeyFinder::init()
 {
 	// Allocate device memory
-	_devCtx = new DeviceContext;
-	_devCtx->init(_device, _numThreads, _numBlocks, _pointsPerThread);
+	_devCtx = new CudaDeviceContext;
+
+	DeviceParameters params;
+	params.device = _device;
+	params.threads = _numThreads;
+	params.blocks = _numBlocks;
+	params.pointsPerThread = _pointsPerThread;
+
+	_devCtx->init(params);
 
 	// Copy points to device
 	generateStartingPoints();
@@ -126,18 +133,18 @@ void KeyFinder::generateStartingPoints()
 	_startingPoints.clear();
 	_iterCount = 0;
 
-	unsigned int totalPoints = _pointsPerThread * _numThreads * _numBlocks;
+	unsigned long long totalPoints = _pointsPerThread * _numThreads * _numBlocks;
 
 	// Generate key pairs for k, k+1, k+2 ... k + <total points in parallel - 1>
 	secp256k1::uint256 privKey = _startExponent;
 
-	for(unsigned int i = 0; i < totalPoints; i++) {
+	for(unsigned long long i = 0; i < totalPoints; i++) {
 		_exponents.push_back(privKey.add(i));
 	}
 
 	secp256k1::generateKeyPairsBulk(secp256k1::G(), _exponents, _startingPoints);
 
-	for(unsigned int i = 0; i < _startingPoints.size(); i++) {
+	for(unsigned long long i = 0; i < _startingPoints.size(); i++) {
 		if(!secp256k1::pointExists(_startingPoints[i])) {
 			throw KeyFinderException("Point generation error");
 		}
@@ -186,6 +193,40 @@ bool KeyFinder::verifyKey(const secp256k1::uint256 &privateKey, const secp256k1:
 	return true;
 }
 
+void KeyFinder::getResults(std::vector<KeyFinderResult> &r)
+{
+	int count = _devCtx->getResultCount();
+
+	if(count == 0) {
+		return;
+	}
+
+	unsigned char *ptr = new unsigned char[count * sizeof(KeyFinderDeviceResult)];
+
+	_devCtx->getResults(ptr, count * sizeof(KeyFinderDeviceResult));
+
+
+	for(int i = 0; i < count; i++) {
+		struct KeyFinderDeviceResult *rPtr = &((struct KeyFinderDeviceResult *)ptr)[i];
+
+		KeyFinderResult minerResult;
+		minerResult.block = rPtr->block;
+		minerResult.thread = rPtr->thread;
+		minerResult.index = rPtr->idx;
+		minerResult.compressed = rPtr->compressed;
+		for(int i = 0; i < 5; i++) {
+			minerResult.hash[i] = rPtr->digest[i];
+		}
+		minerResult.p = secp256k1::ecpoint(secp256k1::uint256(rPtr->x, secp256k1::uint256::BigEndian), secp256k1::uint256(rPtr->y, secp256k1::uint256::BigEndian));
+
+		r.push_back(minerResult);
+	}
+
+	delete[] ptr;
+
+	_devCtx->clearResults();
+}
+
 void KeyFinder::run()
 {
 	unsigned int pointsPerIteration = _numBlocks * _numThreads * _pointsPerThread;
@@ -230,7 +271,7 @@ void KeyFinder::run()
 		if(_devCtx->resultFound()) {
 			std::vector<KeyFinderResult> results;
 
-			_devCtx->getKeyFinderResults(results);
+			getResults(results);
 
 			for(unsigned int i = 0; i < results.size(); i++) {
 				unsigned int index = _devCtx->getIndex(results[i].block, results[i].thread, results[i].index);

diff --git a/KeyFinderLib/KeyFinder.cu b/KeyFinderLib/KeyFinder.cu
@@ -35,8 +35,7 @@ cudaError_t setTargetHash(const unsigned int hash[5])
 	unsigned int h[5];
 
 
-	// Subtract the RIPEmd160 IV from the hash and swp endian, to avoid performing
-	// the addition and endian swap at the end of the RIPEMD160 call
+	// Undo the final round of RIPEMD160 and endian swap to save some computation
 	for(int i = 0; i < 5; i++) {
 		h[i] = swp(hash[i]) - _RIPEMD160_IV_HOST[(i + 1) % 5];
 	}
@@ -88,22 +87,20 @@ __device__ void hashPublicKeyCompressed(const unsigned int *x, const unsigned in
 	ripemd160sha256NoFinal(hash, digestOut);
 }
 
-__device__ void setHashFoundFlag(unsigned int *flagsAra, int idx, int value)
+__device__ void addResult(unsigned int *numResultsPtr, void *results, void *info, int size)
 {
 	grabLock();
-	int threadId = blockIdx.x * blockDim.x + threadIdx.x;
 
-	int base = gridDim.x * blockDim.x * idx;
+	unsigned char *ptr = (unsigned char *)results + (*numResultsPtr);
 
-	flagsAra[base + threadId] = value;
+	memcpy(ptr, info, size);
+
+	(*numResultsPtr)++;
 	releaseLock();
 }
 
-
 __device__ void setResultFound(unsigned int *numResultsPtr, void *results, int idx, bool compressed, unsigned int x[8], unsigned int y[8], unsigned int digest[5])
 {
-	grabLock();
-
 	struct KeyFinderDeviceResult r;
 
 	r.block = blockIdx.x;
@@ -119,11 +116,7 @@ __device__ void setResultFound(unsigned int *numResultsPtr, void *results, int i
 	for(int i = 0; i < 5; i++) {
 		r.digest[i] = endian(digest[i] + _RIPEMD160_IV[(i + 1) % 5]);
 	}
-
-	struct KeyFinderDeviceResult *resultsPtr = (struct KeyFinderDeviceResult *)results;
-	resultsPtr[*numResultsPtr] = r;
-	(*numResultsPtr)++;
-	releaseLock();
+	addResult(numResultsPtr, results, &r, sizeof(r));
 }
 
 __device__ bool checkHash(unsigned int hash[5])
@@ -137,7 +130,7 @@ __device__ bool checkHash(unsigned int hash[5])
 	return equal;
 }
 
-__device__ void doIteration(unsigned int *xPtr, unsigned int *yPtr, unsigned int *chain, int pointsPerThread, unsigned int *numResults, void *results, int flags, int compression)
+__device__ void doIteration(unsigned int *xPtr, unsigned int *yPtr, unsigned int *chain, int pointsPerThread, unsigned int *numResults, void *results, int compression)
 {
 	// Multiply together all (_Gx - x) and then invert
 	unsigned int inverse[8] = { 0,0,0,0,0,0,0,1 };
@@ -149,15 +142,15 @@ __device__ void doIteration(unsigned int *xPtr, unsigned int *yPtr, unsigned int
 		readInt(xPtr, i, x);
 		readInt(yPtr, i, y);
 
-		if(compression == 1 || compression == 2) {
+		if(compression == PointCompressionType::UNCOMPRESSED || compression == PointCompressionType::BOTH) {
 			hashPublicKey(x, y, digest);
 
 			if(checkHash(digest)) {
 				setResultFound(numResults, results, i, false, x, y, digest);
 			}
 		}
 
-		if(compression == 0 || compression == 2) {
+		if(compression == PointCompressionType::COMPRESSED || compression == PointCompressionType::BOTH) {
 			hashPublicKeyCompressed(x, y, digest);
 
 			if(checkHash(digest)) {
@@ -182,7 +175,7 @@ __device__ void doIteration(unsigned int *xPtr, unsigned int *yPtr, unsigned int
 	}
 }
 
-__device__ void doIterationWithDouble(unsigned int *xPtr, unsigned int *yPtr, unsigned int *chain, int pointsPerThread, unsigned int *numResults, void *results, int flags, int compression)
+__device__ void doIterationWithDouble(unsigned int *xPtr, unsigned int *yPtr, unsigned int *chain, int pointsPerThread, unsigned int *numResults, void *results, int compression)
 {
 	// Multiply together all (_Gx - x) and then invert
 	unsigned int inverse[8] = { 0,0,0,0,0,0,0,1 };
@@ -232,12 +225,12 @@ __device__ void doIterationWithDouble(unsigned int *xPtr, unsigned int *yPtr, un
 /**
 * Performs a single iteration
 */
-__global__ void keyFinderKernel(int points, unsigned int flags, unsigned int *x, unsigned int *y, unsigned int *chain, unsigned int *numResults, void *results, int compression)
+__global__ void keyFinderKernel(int points, unsigned int *x, unsigned int *y, unsigned int *chain, unsigned int *numResults, void *results, int compression)
 {
-	doIteration(x, y, chain, points, numResults, results, flags, compression);
+	doIteration(x, y, chain, points, numResults, results, compression);
 }
 
-__global__ void keyFinderKernelWithDouble(int points, unsigned int flags, unsigned int *x, unsigned int *y, unsigned int *chain, unsigned int *numResults, void *results, int compression)
+__global__ void keyFinderKernelWithDouble(int points, unsigned int *x, unsigned int *y, unsigned int *chain, unsigned int *numResults, void *results, int compression)
 {
-	doIterationWithDouble(x, y, chain, points, numResults, results, flags, compression);
+	doIterationWithDouble(x, y, chain, points, numResults, results, compression);
 }
diff --git a/KeyFinderLib/KeyFinder.h b/KeyFinderLib/KeyFinder.h
@@ -5,7 +5,7 @@
 #include "secp256k1.h"
 
 
-class DeviceContext;
+class CudaDeviceContext;
 
 typedef struct {
 	secp256k1::ecpoint publicKey;
@@ -41,6 +41,7 @@ class KeyFinderException {
 	std::string msg;
 };
 
+class KeyFinderResult;
 
 class KeyFinder {
 
@@ -56,7 +57,7 @@ class KeyFinder {
 
 	unsigned int _statusInterval;
 
-	DeviceContext *_devCtx;
+	CudaDeviceContext *_devCtx;
 
 	unsigned long long _iterCount;
 	unsigned long long _total;
@@ -94,6 +95,7 @@ class KeyFinder {
 
 	bool verifyKey(const secp256k1::uint256 &privateKey, const secp256k1::ecpoint &publicKey, const unsigned int hash[5], bool compressed);
 
+	void getResults(std::vector<KeyFinderResult> &r);
 
 public:
 	class Compression {

diff --git a/KeyFinderLib/cudabridge.cu b/KeyFinderLib/cudabridge.cu
@@ -1,15 +1,15 @@
 #include "cudabridge.h"
 
 
-__global__ void keyFinderKernel(int points, unsigned int flags, unsigned int *x, unsigned int *y, unsigned int *chain, unsigned int *numResults, void *results, int compression);
-__global__ void keyFinderKernelWithDouble(int points, unsigned int flags, unsigned int *x, unsigned int *y, unsigned int *chain, unsigned int *numResults, void *results, int compression);
+__global__ void keyFinderKernel(int points, unsigned int *x, unsigned int *y, unsigned int *chain, unsigned int *numResults, void *results, int compression);
+__global__ void keyFinderKernelWithDouble(int points, unsigned int *x, unsigned int *y, unsigned int *chain, unsigned int *numResults, void *results, int compression);
 
 void callKeyFinderKernel(KernelParams &params, bool useDouble, int compression)
 {
 	if(useDouble) {
-		keyFinderKernelWithDouble <<<params.blocks, params.threads >> >(params.points, params.flags, params.x, params.y, params.chain, params.numResults, params.results, compression);
+		keyFinderKernelWithDouble <<<params.blocks, params.threads >> >(params.points, params.x, params.y, params.chain, params.numResults, params.results, compression);
 	} else {
-		keyFinderKernel <<<params.blocks, params.threads >> > (params.points, params.flags, params.x, params.y, params.chain, params.numResults, params.results, compression);
+		keyFinderKernel <<<params.blocks, params.threads >> > (params.points, params.x, params.y, params.chain, params.numResults, params.results, compression);
 	}
 	waitForKernel();
 }