Initialize on GPU

-Initialize the points on the GPU instead of CPU for faster startup
-Fixed bug where device detection could fail

KeyFinder/main.cpp

@@ -164,7 +164,8 @@ int main(int argc, char **argv)
 	secp256k1::uint256 start(1);
 	unsigned long long range = 0;
 
-	if(cuda::getDeviceCount == 0) {
+
+	if(cuda::getDeviceCount() == 0) {
 		Logger::log(LogLevel::Error, "No CUDA devices available");
 		return 1;
 	}
@@ -305,7 +306,7 @@ int main(int argc, char **argv)
 		f.init();
 		f.run();
 	} catch(KeyFinderException ex) {
-		Logger::log(LogLevel::Info, "Exiting");
+		Logger::log(LogLevel::Info, "Error: " + ex.msg + " Exiting.");
 		return 1;
 	}
 

KeyFinderLib/KeyFinder.cpp

@@ -17,6 +17,7 @@
 
 #include "Logger.h"
 
+#include "ec.h"
 
 void KeyFinder::defaultResultCallback(KeyFinderResultInfo result)
 {
@@ -191,8 +192,6 @@ void KeyFinder::init()
 	// Copy points to device
 	generateStartingPoints();
 
-	_devCtx->copyPoints(_startingPoints);
-
 	setTargetsOnDevice();
 
 	allocateChainBuf(_numThreads * _numBlocks * _pointsPerThread);
@@ -228,7 +227,7 @@ void KeyFinder::generateStartingPoints()
 	unsigned long long totalPoints = _pointsPerThread * _numThreads * _numBlocks;
 	unsigned long long totalMemory = totalPoints * 40;
 
-	Logger::log(LogLevel::Info, "Generating " + util::formatThousands(totalPoints) + " starting points (" + util::format("%.1f", (double)totalMemory/(double)(1024*1024)) + "MB)");
+	Logger::log(LogLevel::Info, "Generating " + util::formatThousands(totalPoints) + " starting points (" + util::format("%.1f", (double)totalMemory / (double)(1024 * 1024)) + "MB)");
 
 	// Generate key pairs for k, k+1, k+2 ... k + <total points in parallel - 1>
 	secp256k1::uint256 privKey = _startExponent;
@@ -237,13 +236,31 @@ void KeyFinder::generateStartingPoints()
 		_exponents.push_back(privKey.add(i));
 	}
 
-	secp256k1::generateKeyPairsBulk(secp256k1::G(), _exponents, _startingPoints);
+	_deviceKeys.init(_numBlocks, _numThreads, _pointsPerThread, _exponents);
 
-	for(unsigned int i = 0; i < _startingPoints.size(); i++) {
-		if(!secp256k1::pointExists(_startingPoints[i])) {
-			throw KeyFinderException("Point generation error");
+	// Show progress in 10% increments
+	double pct = 10.0;
+	for(int i = 1; i <= 256; i++) {
+		_deviceKeys.doStep();
+		if(((double)i / 256.0) * 100.0 >= pct) {
+			Logger::log(LogLevel::Info, util::format("%.1f%%", pct));
+			pct += 10.0;
 		}
 	}
+
+#ifdef _DEBUG
+	try {
+		Logger::log(LogLevel::Debug, "Verifying points on device...");
+		_deviceKeys.selfTest(_exponents);
+	} catch(std::string &e) {
+		Logger::log(LogLevel::Debug, e);
+		exit(1);
+	}
+#endif
+
+	Logger::log(LogLevel::Info, "Done");
+
+	_deviceKeys.clearPrivateKeys();
 }
 
 
@@ -311,7 +328,6 @@ void KeyFinder::getResults(std::vector<KeyFinderResult> &r)
 
 	unsigned char *ptr = new unsigned char[count * sizeof(KeyFinderDeviceResult)];
 
-	//_devCtx->getResults(ptr, count * sizeof(KeyFinderDeviceResult));
 	_resultList.read(ptr, count);
 
 	for(int i = 0; i < count; i++) {
@@ -443,8 +459,10 @@ void KeyFinder::run()
 				}
 
 				// Update hash targets on device
-				Logger::log(LogLevel::Info, "Reloading targets");
-				setTargetsOnDevice();
+				if(_targets.size() > 0) {
+					Logger::log(LogLevel::Info, "Reloading targets");
+					setTargetsOnDevice();
+				}
 			}
 		}
 		_iterCount++;

KeyFinderLib/KeyFinder.cu

@@ -15,6 +15,7 @@
 
 #include "hashlookup.cuh"
 #include "atomiclist.cuh"
+#include "ec.cuh"
 
 __constant__ unsigned int _INC_X[8];
 
@@ -138,9 +139,11 @@ __device__ void setResultFound(int idx, bool compressed, unsigned int x[8], unsi
 	atomicListAdd(&r, sizeof(r));
 }
 
-__device__ void doIteration(unsigned int *xPtr, unsigned int *yPtr, int pointsPerThread, int compression)
+__device__ void doIteration(int pointsPerThread, int compression)
 {
 	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 };
@@ -172,7 +175,7 @@ __device__ void doIteration(unsigned int *xPtr, unsigned int *yPtr, int pointsPe
 			}
 		}
 
-		beginBatchAdd(_INC_X, x, chain, i, inverse);
+		beginBatchAdd(_INC_X, x, chain, i, i, inverse);
 	}
 
 	doBatchInverse(inverse);
@@ -182,16 +185,18 @@ __device__ void doIteration(unsigned int *xPtr, unsigned int *yPtr, int pointsPe
 		unsigned int newX[8];
 		unsigned int newY[8];
 
-		completeBatchAdd(_INC_X, _INC_Y, xPtr, yPtr, i, chain, inverse, newX, newY);
+		completeBatchAdd(_INC_X, _INC_Y, xPtr, yPtr, i, i, chain, inverse, newX, newY);
 
 		writeInt(xPtr, i, newX);
 		writeInt(yPtr, i, newY);
 	}
 }
 
-__device__ void doIterationWithDouble(unsigned int *xPtr, unsigned int *yPtr, int pointsPerThread, int compression)
+__device__ void doIterationWithDouble(int pointsPerThread, int compression)
 {
 	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 };
@@ -201,7 +206,6 @@ __device__ void doIterationWithDouble(unsigned int *xPtr, unsigned int *yPtr, in
 		unsigned int digest[5];
 
 		readInt(xPtr, i, x);
-
 
 		// uncompressed
 		if(compression == PointCompressionType::UNCOMPRESSED || compression == PointCompressionType::BOTH) {
@@ -216,17 +220,19 @@ __device__ void doIterationWithDouble(unsigned int *xPtr, unsigned int *yPtr, in
 
 		// 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, inverse);
+		beginBatchAddWithDouble(_INC_X, _INC_Y, xPtr, chain, i, i, inverse);
 	}
 
 	doBatchInverse(inverse);
@@ -236,7 +242,7 @@ __device__ void doIterationWithDouble(unsigned int *xPtr, unsigned int *yPtr, in
 		unsigned int newX[8];
 		unsigned int newY[8];
 
-		completeBatchAddWithDouble(_INC_X, _INC_Y, xPtr, yPtr, i, chain, inverse, newX, newY);
+		completeBatchAddWithDouble(_INC_X, _INC_Y, xPtr, yPtr, i, i, chain, inverse, newX, newY);
 
 		writeInt(xPtr, i, newX);
 		writeInt(yPtr, i, newY);
@@ -248,10 +254,10 @@ __device__ void doIterationWithDouble(unsigned int *xPtr, unsigned int *yPtr, in
 */
 __global__ void keyFinderKernel(int points, unsigned int *x, unsigned int *y, int compression)
 {
-	doIteration(x, y, points, compression);
+	doIteration(points, compression);
 }
 
 __global__ void keyFinderKernelWithDouble(int points, unsigned int *x, unsigned int *y, int compression)
 {
-	doIterationWithDouble(x, y, points, compression);
+	doIterationWithDouble(points, compression);
 }

KeyFinderLib/KeyFinder.h

@@ -6,6 +6,7 @@
 #include "secp256k1.h"
 #include "atomiclist.h"
 #include "hashlookup.h"
+#include "ec.h"
 
 class CudaDeviceContext;
 
@@ -117,6 +118,8 @@ class KeyFinder {
 
 private:
 
+	CudaDeviceKeys _deviceKeys;
+
 	CudaAtomicList _resultList;
 
 	CudaHashLookup _targetLookup;

KeyFinderLib/KeyFinderLib.vcxproj

@@ -21,6 +21,8 @@
   <ItemGroup>
     <ClInclude Include="atomiclist.h" />
     <ClInclude Include="cudabridge.h" />
+    <ClInclude Include="ec.cuh" />
+    <ClInclude Include="ec.h" />
     <ClInclude Include="hashlookup.h" />
     <ClInclude Include="KeyFinder.h" />
     <ClInclude Include="hashlookup.cuh" />
@@ -49,6 +51,7 @@
   </ItemGroup>
   <ItemGroup>
     <CudaCompile Include="cudabridge.cu" />
+    <CudaCompile Include="ec.cu" />
     <CudaCompile Include="KeyFinder.cu" />
     <CudaCompile Include="hashlookup.cu" />
     <CudaCompile Include="atomiclist.cu" />