Add support for EthereumStratum/1.0.0 mode (see -ES option). This adds compatibility with NiceHash.com.

Author: kenshirothefist

cmake/scripts/macdeployfix.sh

@@ -201,6 +201,10 @@ class MinerCLI
 				BOOST_THROW_EXCEPTION(BadArgument());
 			}
 		}
+		else if (arg == "-ES" || arg == "--ethereum-stratum")
+		{
+			m_ethereumStratum = true;
+		}
 		else if ((arg == "-FO" || arg == "--failover-userpass") && i + 1 < argc)
 		{
 			string userpass = string(argv[++i]);
@@ -330,6 +334,11 @@ class MinerCLI
 			string mode = argv[++i];
 			if (mode == "parallel") m_dagLoadMode = DAG_LOAD_MODE_PARALLEL;
 			else if (mode == "sequential") m_dagLoadMode = DAG_LOAD_MODE_SEQUENTIAL;
+			else if (mode == "single")
+			{
+				m_dagLoadMode = DAG_LOAD_MODE_SINGLE;
+				m_dagCreateDevice = stol(argv[++i]);
+			}
 			else
 			{
 				cerr << "Bad " << arg << " option: " << argv[i] << endl;
@@ -531,7 +540,8 @@ class MinerCLI
 					m_clAllowCPU,
 					m_extraGPUMemory,
 					0,
-					m_dagLoadMode
+					m_dagLoadMode,
+					m_dagCreateDevice
 				))
 				exit(1);
 			EthashGPUMiner::setNumInstances(m_miningThreads);
@@ -557,7 +567,8 @@ class MinerCLI
 				m_extraGPUMemory,
 				m_cudaSchedule,
 				0,
-				m_dagLoadMode
+				m_dagLoadMode,
+				m_dagCreateDevice
 				))
 				exit(1);
 #else
@@ -592,7 +603,8 @@ class MinerCLI
 			<< "    -O, --userpass <username.workername:password> Stratum login credentials" << endl
 			<< "    -FO, --failover-userpass <username.workername:password> Failover stratum login credentials (optional, will use normal credentials when omitted)" << endl
 			<< "    --work-timeout <n> reconnect/failover after n seconds of working on the same (stratum) job. Defaults to 180. Don't set lower than max. avg. block time" << endl
-			<< "    -SV, --stratum-version <n>  Stratum client version. Defaults to 1 (async client). Use 2 to test new synchronous client."
+			<< "    -SV, --stratum-version <n>  Stratum client version. Defaults to 1 (async client). Use 2 to test new synchronous client." << endl
+			<< "    -ES, --ethereum-stratum  Use EthereumStratum/1.0.0 mode." << endl
 #endif
 #if ETH_JSONRPC || ETH_STRATUM || !ETH_TRUE
 			<< "    --farm-recheck <n>  Leave n ms between checks for changed work (default: 500). When using stratum, use a high value (i.e. 2000) to get more stable hashrate output" << endl
@@ -618,6 +630,7 @@ class MinerCLI
 			<< "    -L, --dag-load-mode <mode> DAG generation mode." << endl
 			<< "        parallel    - load DAG on all GPUs at the same time (default)" << endl
 			<< "        sequential  - load DAG on GPUs one after another. Use this when the miner crashes during DAG generation" << endl
+			<< "        single <n>  - generate DAG on device n, then copy to other devices" << endl
 #if ETH_ETHASHCL || !ETH_TRUE
 			<< "    --cl-extragpu-mem Set the memory (in MB) you believe your GPU requires for stuff other than mining. default: 0" << endl
 			<< "    --cl-local-work Set the OpenCL local work size. Default is " << toString(ethash_cl_miner::c_defaultLocalWorkSize) << endl
@@ -992,7 +1005,7 @@ class MinerCLI
 
 		// this is very ugly, but if Stratum Client V2 tunrs out to be a success, V1 will be completely removed anyway
 		if (m_stratumClientVersion == 1) {
-			EthStratumClient client(&f, m_minerType, m_farmURL, m_port, m_user, m_pass, m_maxFarmRetries, m_worktimeout);
+			EthStratumClient client(&f, m_minerType, m_farmURL, m_port, m_user, m_pass, m_maxFarmRetries, m_worktimeout, m_ethereumStratum);
 			if (m_farmFailOverURL != "")
 			{
 				if (m_fuser != "")
@@ -1032,7 +1045,7 @@ class MinerCLI
 			}
 		}
 		else if (m_stratumClientVersion == 2) {
-			EthStratumClientV2 client(&f, m_minerType, m_farmURL, m_port, m_user, m_pass, m_maxFarmRetries, m_worktimeout);
+			EthStratumClientV2 client(&f, m_minerType, m_farmURL, m_port, m_user, m_pass, m_maxFarmRetries, m_worktimeout, m_ethereumStratum);
 			if (m_farmFailOverURL != "")
 			{
 				if (m_fuser != "")
@@ -1100,6 +1113,7 @@ class MinerCLI
 	// default value was 350MB of GPU memory for other stuff (windows system rendering, e.t.c.)
 	unsigned m_extraGPUMemory = 0;// 350000000; don't assume miners run desktops...
 	unsigned m_dagLoadMode = 0; // parallel
+	unsigned m_dagCreateDevice = 0;
 	/// Benchmarking params
 	bool m_phoneHome = false;
 	unsigned m_benchmarkWarmup = 15;
@@ -1121,6 +1135,7 @@ class MinerCLI
 
 #if ETH_STRATUM || !ETH_TRUE
 	int m_stratumClientVersion = 1;
+	bool m_ethereumStratum = false;
 	string m_user;
 	string m_pass;
 	string m_port;

ethminer/MinerAux.h

@@ -36,3 +36,4 @@ cd ..
 
 goto :EOF
 
+

extdep/getstuff.bat

@@ -504,7 +504,7 @@ typedef struct
 	unsigned buf;
 } pending_batch;
 
-void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook& hook)
+void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook& hook, bool _ethStratum, uint64_t _startN)
 {
 	try
 	{
@@ -531,7 +531,9 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook
 
 		unsigned buf = 0;
 		random_device engine;
-		uint64_t start_nonce = uniform_int_distribution<uint64_t>()(engine);
+		uint64_t start_nonce;
+		if (_ethStratum) start_nonce = _startN;
+		else start_nonce = uniform_int_distribution<uint64_t>()(engine);
 		for (;; start_nonce += m_globalWorkSize)
 		{
 			// supply output buffer to kernel

libethash-cl/ethash_cl_miner.cpp

@@ -59,7 +59,7 @@ class ethash_cl_miner
 		unsigned _deviceId
 		);
 	void finish();
-	void search(uint8_t const* _header, uint64_t _target, search_hook& _hook);
+	void search(uint8_t const* _header, uint64_t _target, search_hook& _hook, bool _ethStratum, uint64_t _startN);
 
 	/* -- default values -- */
 	/// Default value of the local work size. Also known as workgroup size.

libethash-cl/ethash_cl_miner.h

@@ -199,7 +199,7 @@ void ethash_cuda_miner::finish()
 	CUDA_SAFE_CALL(cudaDeviceReset());
 }
 
-bool ethash_cuda_miner::init(ethash_light_t _light, uint8_t const* _lightData, uint64_t _lightSize, unsigned _deviceId)
+bool ethash_cuda_miner::init(ethash_light_t _light, uint8_t const* _lightData, uint64_t _lightSize, unsigned _deviceId, bool _cpyToHost, volatile void** hostDAG)
 {
 	try
 	{
@@ -229,10 +229,14 @@ bool ethash_cuda_miner::init(ethash_light_t _light, uint8_t const* _lightData, u
 		uint32_t lightSize64 = (unsigned)(_lightSize / sizeof(node));
 
 		// create buffer for cache
-		hash64_t * light;
-		CUDA_SAFE_CALL(cudaMalloc(reinterpret_cast<void**>(&light), _lightSize));
-		// copy dag cache to CPU.
-		CUDA_SAFE_CALL(cudaMemcpy(reinterpret_cast<void*>(light), _lightData, _lightSize, cudaMemcpyHostToDevice));
+		hash64_t * light = NULL;
+
+		if (!*hostDAG)
+		{
+			CUDA_SAFE_CALL(cudaMalloc(reinterpret_cast<void**>(&light), _lightSize));
+			// copy dag cache to CPU.
+			CUDA_SAFE_CALL(cudaMemcpy(reinterpret_cast<void*>(light), _lightData, _lightSize, cudaMemcpyHostToDevice));
+		}
 
 		// create buffer for dag
 		hash128_t * dag;
@@ -252,9 +256,28 @@ bool ethash_cuda_miner::init(ethash_light_t _light, uint8_t const* _lightData, u
 
 		m_sharedBytes = device_props.major * 100 < SHUFFLE_MIN_VER ? (64 * s_blockSize) / 8 : 0 ;
 
+		if (!*hostDAG)
+		{
+			cout << "Generating DAG for GPU #" << device_num << endl;
+			ethash_generate_dag(dagSize, s_gridSize, s_blockSize, m_streams[0], device_num);
+
+			if (_cpyToHost)
+			{
+				uint8_t* memoryDAG = new uint8_t[dagSize];
+				if (!memoryDAG) throw std::runtime_error("Failed to init host memory for DAG, not enough memory?");
+
+				cout << "Copying DAG from GPU #" << device_num << " to host" << endl;
+				CUDA_SAFE_CALL(cudaMemcpy(reinterpret_cast<void*>(memoryDAG), dag, dagSize, cudaMemcpyDeviceToHost));
 
-		cout << "Generating DAG for GPU #" << device_num << endl;
-		ethash_generate_dag(dagSize, s_gridSize, s_blockSize, m_streams[0], device_num);
+				*hostDAG = (void*)memoryDAG;
+			}
+		}
+		else
+		{
+			cout << "Copying DAG from host to GPU #" << device_num << endl;
+			const void* hdag = (const void*)(*hostDAG);
+			CUDA_SAFE_CALL(cudaMemcpy(reinterpret_cast<void*>(dag), hdag, dagSize, cudaMemcpyHostToDevice));
+		}
 
 		return true;
 	}
@@ -264,7 +287,7 @@ bool ethash_cuda_miner::init(ethash_light_t _light, uint8_t const* _lightData, u
 	}
 }
 
-void ethash_cuda_miner::search(uint8_t const* header, uint64_t target, search_hook& hook)
+void ethash_cuda_miner::search(uint8_t const* header, uint64_t target, search_hook& hook, bool _ethStratum, uint64_t _startN)
 {
 	bool initialize = false;
 	bool exit = false;
@@ -280,14 +303,34 @@ void ethash_cuda_miner::search(uint8_t const* header, uint64_t target, search_ho
 		set_target(m_current_target);
 		initialize = true;
 	}
-	if (initialize)
+	if (_ethStratum)
 	{
-		random_device engine;
-		m_current_nonce = uniform_int_distribution<uint64_t>()(engine);
-		m_current_index = 0;
-		CUDA_SAFE_CALL(cudaDeviceSynchronize());
-		for (unsigned int i = 0; i < s_numStreams; i++)
-			m_search_buf[i][0] = 0;
+		if (initialize)
+		{
+			m_starting_nonce = 0;
+			m_current_index = 0;
+			CUDA_SAFE_CALL(cudaDeviceSynchronize());
+			for (unsigned int i = 0; i < s_numStreams; i++)
+				m_search_buf[i][0] = 0;
+		}
+		if (m_starting_nonce != _startN)
+		{
+			// reset nonce counter
+			m_starting_nonce = _startN;
+			m_current_nonce = m_starting_nonce;
+		}
+	}
+	else
+	{
+		if (initialize)
+		{
+			random_device engine;
+			m_current_nonce = uniform_int_distribution<uint64_t>()(engine);
+			m_current_index = 0;
+			CUDA_SAFE_CALL(cudaDeviceSynchronize());
+			for (unsigned int i = 0; i < s_numStreams; i++)
+				m_search_buf[i][0] = 0;
+		}
 	}
 	uint64_t batch_size = s_gridSize * s_blockSize;
 	for (; !exit; m_current_index++, m_current_nonce += batch_size)

libethash-cuda/ethash_cuda_miner.cpp

@@ -35,10 +35,10 @@ class ethash_cuda_miner
 		uint64_t _currentBlock
 		);
 
-	bool init(ethash_light_t _light, uint8_t const* _lightData, uint64_t _lightSize, unsigned _deviceId);
+	bool init(ethash_light_t _light, uint8_t const* _lightData, uint64_t _lightSize, unsigned _deviceId, bool _cpyToHost, volatile void** hostDAG);
 
 	void finish();
-	void search(uint8_t const* header, uint64_t target, search_hook& hook);
+	void search(uint8_t const* header, uint64_t target, search_hook& hook, bool _ethStratum, uint64_t _startN);
 
 	/* -- default values -- */
 	/// Default value of the block size. Also known as workgroup size.
@@ -52,6 +52,7 @@ class ethash_cuda_miner
 	hash32_t m_current_header;
 	uint64_t m_current_target;
 	uint64_t m_current_nonce;
+	uint64_t m_starting_nonce;
 	uint64_t m_current_index;
 
 	uint32_t m_sharedBytes;

libethash-cuda/ethash_cuda_miner.h

@@ -53,17 +53,20 @@ struct EthashProofOfWork
 	struct WorkPackage
 	{
 		WorkPackage() = default;
-		WorkPackage(Ethash::BlockHeader const& _bh):
+		WorkPackage(Ethash::BlockHeader const& _bh) :
 			boundary(_bh.boundary()),
 			headerHash(_bh.hashWithout()),
 			seedHash(_bh.seedHash())
-		{}
+		{ }
 		void reset() { headerHash = h256(); }
 		operator bool() const { return headerHash != h256(); }
 
 		h256 boundary;
 		h256 headerHash;	///< When h256() means "pause until notified a new work package is available".
 		h256 seedHash;
+
+		uint64_t startNonce = 0;
+		int exSizeBits = -1;
 	};
 
 	static const WorkPackage NullWorkPackage;

libethcore/EthashAux.h

@@ -144,32 +144,62 @@ void EthashCUDAMiner::workLoop()
 		cnote << "set work; seed: " << "#" + w.seedHash.hex().substr(0, 8) + ", target: " << "#" + w.boundary.hex().substr(0, 16);
 		if (!m_miner || m_minerSeed != w.seedHash)
 		{
+			unsigned device = s_devices[index()] > -1 ? s_devices[index()] : index();
+
 			if (s_dagLoadMode == DAG_LOAD_MODE_SEQUENTIAL)
 			{
 				while (s_dagLoadIndex < index()) {
 					this_thread::sleep_for(chrono::seconds(1));
 				}
 			}
+			else if (s_dagLoadMode == DAG_LOAD_MODE_SINGLE)
+			{
+				if (device != s_dagCreateDevice)
+				{
+					// wait until DAG is created on selected device
+					while (s_dagInHostMemory == NULL) {
+						this_thread::sleep_for(chrono::seconds(1));
+					}
+				}
+				else
+				{
+					// reset load index
+					s_dagLoadIndex = 0;
+				}
+			}
 
 			cnote << "Initialising miner...";
 			m_minerSeed = w.seedHash;
 
 			delete m_miner;
 			m_miner = new ethash_cuda_miner;
 
-			unsigned device = s_devices[index()] > -1 ? s_devices[index()] : index();
-
 			EthashAux::LightType light;
 			light = EthashAux::light(w.seedHash);
 			//bytesConstRef dagData = dag->data();
 			bytesConstRef lightData = light->data();
 
-			m_miner->init(light->light, lightData.data(), lightData.size(), device);
+			m_miner->init(light->light, lightData.data(), lightData.size(), device, (s_dagLoadMode == DAG_LOAD_MODE_SINGLE), &s_dagInHostMemory);
 			s_dagLoadIndex++;
+
+			if (s_dagLoadMode == DAG_LOAD_MODE_SINGLE)
+			{
+				if (s_dagLoadIndex >= s_numInstances && s_dagInHostMemory)
+				{
+					// all devices have loaded DAG, we can free now
+					delete[] s_dagInHostMemory;
+					s_dagInHostMemory = NULL;
+
+					cout << "Freeing DAG from host" << endl;
+				}
+			}
 		}
 
 		uint64_t upper64OfBoundary = (uint64_t)(u64)((u256)w.boundary >> 192);
-		m_miner->search(w.headerHash.data(), upper64OfBoundary, *m_hook);
+		uint64_t startN;
+		if (w.exSizeBits >= 0)
+			startN = w.startNonce | ((uint64_t)index() << (64 - 4 - w.exSizeBits)); // this can support up to 16 devices
+		m_miner->search(w.headerHash.data(), upper64OfBoundary, *m_hook, (w.exSizeBits >= 0), startN);
 	}
 	catch (std::runtime_error const& _e)
 	{
@@ -207,10 +237,12 @@ bool EthashCUDAMiner::configureGPU(
 	unsigned _extraGPUMemory,
 	unsigned _scheduleFlag,
 	uint64_t _currentBlock,
-	unsigned _dagLoadMode
+	unsigned _dagLoadMode,
+	unsigned _dagCreateDevice
 	)
 {
 	s_dagLoadMode = _dagLoadMode;
+	s_dagCreateDevice = _dagCreateDevice;
 	_blockSize = ((_blockSize + 7) / 8) * 8;
 
 	if (!ethash_cuda_miner::configureGPU(

libethcore/EthashCUDAMiner.cpp

@@ -54,7 +54,8 @@ namespace eth
 			unsigned _extraGPUMemory,
 			unsigned _scheduleFlag,
 			uint64_t _currentBlock,
-			unsigned _dagLoadMode
+			unsigned _dagLoadMode,
+			unsigned _dagCreateDevice
 			);
 		static void setNumInstances(unsigned _instances) 
 		{