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

ethminer/MinerAux.h

@@ -201,6 +201,10 @@ public:
 				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 @@ public:
 			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 @@ public:
 					m_clAllowCPU,
 					m_extraGPUMemory,
 					0,
-					m_dagLoadMode
+					m_dagLoadMode,
+					m_dagCreateDevice
 				))
 				exit(1);
 			EthashGPUMiner::setNumInstances(m_miningThreads);
@@ -557,7 +567,8 @@ public:
 				m_extraGPUMemory,
 				m_cudaSchedule,
 				0,
-				m_dagLoadMode
+				m_dagLoadMode,
+				m_dagCreateDevice
 				))
 				exit(1);
 #else
@@ -592,7 +603,8 @@ public:
 			<< "    -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 @@ public:
 			<< "    -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 @@ private:
 
 		// 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 @@ private:
 			}
 		}
 		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 @@ private:
 	// 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 @@ private:
 
 #if ETH_STRATUM || !ETH_TRUE
 	int m_stratumClientVersion = 1;
+	bool m_ethereumStratum = false;
 	string m_user;
 	string m_pass;
 	string m_port;

extdep/getstuff.bat

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

libethash-cl/ethash_cl_miner.cpp

@@ -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.h

@@ -59,7 +59,7 @@ public:
 		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-cuda/ethash_cuda_miner.cpp

@@ -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.h

@@ -35,10 +35,10 @@ public:
 		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 @@ private:
 	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;

libethcore/EthashAux.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/EthashCUDAMiner.cpp

@@ -144,12 +144,29 @@ 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;
@@ -157,19 +174,32 @@ void EthashCUDAMiner::workLoop()
 			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.h

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

libethcore/EthashGPUMiner.cpp

@@ -183,7 +183,10 @@ void EthashGPUMiner::workLoop()
 		}
 
 		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 (cl::Error const& _e)
 	{
@@ -222,10 +225,12 @@ bool EthashGPUMiner::configureGPU(
 	bool _allowCPU,
 	unsigned _extraGPUMemory,
 	uint64_t _currentBlock,
-	unsigned _dagLoadMode
+	unsigned _dagLoadMode,
+	unsigned _dagCreateDevice
 )
 {
 	s_dagLoadMode = _dagLoadMode;
+	s_dagCreateDevice = _dagCreateDevice;
 
 	s_platformId = _platformId;
 	s_deviceId = _deviceId;

libethcore/EthashGPUMiner.h

@@ -53,7 +53,8 @@ public:
 		bool _allowCPU,
 		unsigned _extraGPUMemory,
 		uint64_t _currentBlock,
-		unsigned _dagLoadMode
+		unsigned _dagLoadMode,
+		unsigned _dagCreateDevice
 	);
 	static void setNumInstances(unsigned _instances) { s_numInstances = std::min<unsigned>(_instances, getNumDevices()); }
 	static void setDevices(unsigned * _devices, unsigned _selectedDeviceCount)

libethcore/Farm.h

@@ -67,7 +67,7 @@ public:
 	void setWork(WorkPackage const& _wp)
 	{
 		WriteGuard l(x_minerWork);
-		if (_wp.headerHash == m_work.headerHash)
+		if (_wp.headerHash == m_work.headerHash && _wp.startNonce == m_work.startNonce)
 			return;
 		m_work = _wp;
 		for (auto const& m: m_miners)

libethcore/Miner.cpp

@@ -8,6 +8,12 @@ template <>
 unsigned dev::eth::GenericMiner<dev::eth::EthashProofOfWork>::s_dagLoadMode = 0;
 
 template <>
-unsigned dev::eth::GenericMiner<dev::eth::EthashProofOfWork>::s_dagLoadIndex = 0;
+volatile unsigned dev::eth::GenericMiner<dev::eth::EthashProofOfWork>::s_dagLoadIndex = 0;
+
+template <>
+unsigned dev::eth::GenericMiner<dev::eth::EthashProofOfWork>::s_dagCreateDevice = 0;
+
+template <>
+volatile void* dev::eth::GenericMiner<dev::eth::EthashProofOfWork>::s_dagInHostMemory = NULL;
 
 

libethcore/Miner.h

@@ -38,6 +38,7 @@
 
 #define DAG_LOAD_MODE_PARALLEL	 0
 #define DAG_LOAD_MODE_SEQUENTIAL 1
+#define DAG_LOAD_MODE_SINGLE	 2
 
 
 using namespace std;
@@ -218,7 +219,9 @@ protected:
 	void accumulateHashes(unsigned _n) { m_hashCount += _n; }
 
 	static unsigned s_dagLoadMode;
-	static unsigned s_dagLoadIndex;
+	static volatile unsigned s_dagLoadIndex;
+	static unsigned s_dagCreateDevice;
+	static volatile void* s_dagInHostMemory;
 private:
 	FarmFace* m_farm = nullptr;
 	unsigned m_index;

libstratum/EthStratumClient.cpp

@@ -1,10 +1,33 @@
  
 #include "EthStratumClient.h"
 #include <libdevcore/Log.h>
+#include <libethash/endian.h>
 using boost::asio::ip::tcp;
 
 
-EthStratumClient::EthStratumClient(GenericFarm<EthashProofOfWork> * f, MinerType m, string const & host, string const & port, string const & user, string const & pass, int const & retries, int const & worktimeout)
+static void diffToTarget(uint32_t *target, double diff)
+{
+	uint32_t target2[8];
+	uint64_t m;
+	int k;
+
+	for (k = 6; k > 0 && diff > 1.0; k--)
+		diff /= 4294967296.0;
+	m = (uint64_t)(4294901760.0 / diff);
+	if (m == 0 && k == 6)
+		memset(target2, 0xff, 32);
+	else {
+		memset(target2, 0, 32);
+		target2[k] = (uint32_t)m;
+		target2[k + 1] = (uint32_t)(m >> 32);
+	}
+
+	for (int i = 0; i < 32; i++)
+		((uint8_t*)target)[31 - i] = ((uint8_t*)target2)[i];
+}
+
+
+EthStratumClient::EthStratumClient(GenericFarm<EthashProofOfWork> * f, MinerType m, string const & host, string const & port, string const & user, string const & pass, int const & retries, int const & worktimeout, bool ethstratum)
 	: m_socket(m_io_service)
 {
 	m_minerType = m;
@@ -21,6 +44,8 @@ EthStratumClient::EthStratumClient(GenericFarm<EthashProofOfWork> * f, MinerType
 	m_maxRetries = retries;
 	m_worktimeout = worktimeout;
 
+	m_ethereumStratum = ethstratum;
+
 	p_farm = f;
 	p_worktimer = nullptr;
 	connect();
@@ -154,7 +179,10 @@ void EthStratumClient::connect_handler(const boost::system::error_code& ec, tcp:
 			}
 		}
 		std::ostream os(&m_requestBuffer);
-		os << "{\"id\": 1, \"method\": \"mining.subscribe\", \"params\": []}\n";
+		if (m_ethereumStratum)
+			os << "{\"id\": 1, \"method\": \"mining.subscribe\", \"params\": [\"ethminer/" << ETH_PROJECT_VERSION << "\",\"EthereumStratum/1.0.0\"]}\n";
+		else
+			os << "{\"id\": 1, \"method\": \"mining.subscribe\", \"params\": []}\n";
 
 		
 		async_write(m_socket, m_requestBuffer,
@@ -236,6 +264,16 @@ void EthStratumClient::readResponse(const boost::system::error_code& ec, std::si
 	}
 }
 
+void EthStratumClient::processExtranonce(std::string& enonce)
+{
+	m_extraNonceHexSize = enonce.length();
+
+	cnote << "Extranonce set to " << enonce;
+
+	for (int i = enonce.length(); i < 16; ++i) enonce += "0";
+	m_extraNonce = h64(enonce);
+}
+
 void EthStratumClient::processReponse(Json::Value& responseObject)
 {
 	Json::Value error = responseObject.get("error", new Json::Value);
@@ -252,13 +290,29 @@ void EthStratumClient::processReponse(Json::Value& responseObject)
 	case 1:
 		cnote << "Subscribed to stratum server";
 
-		os << "{\"id\": 2, \"method\": \"mining.authorize\", \"params\": [\"" << p_active->user << "\",\"" << p_active->pass << "\"]}\n";
+		if (m_ethereumStratum)
+		{
+			m_nextWorkDifficulty = 1;
+			params = responseObject.get("result", Json::Value::null);
+			if (params.isArray())
+			{
+				std::string enonce = params.get((Json::Value::ArrayIndex)1, "").asString();
+				processExtranonce(enonce);
+			}
+
+			os << "{\"id\": 2, \"method\": \"mining.extranonce.subscribe\", \"params\": []}\n";
+		}
+
+		os << "{\"id\": 3, \"method\": \"mining.authorize\", \"params\": [\"" << p_active->user << "\",\"" << p_active->pass << "\"]}\n";
 
 		async_write(m_socket, m_requestBuffer,
 			boost::bind(&EthStratumClient::handleResponse, this,
 			boost::asio::placeholders::error));
 		break;
 	case 2:
+		// nothing to do...
+		break;
+	case 3:
 		m_authorized = responseObject.get("result", Json::Value::null).asBool();
 		if (!m_authorized)
 		{
@@ -286,55 +340,107 @@ void EthStratumClient::processReponse(Json::Value& responseObject)
 			if (params.isArray())
 			{
 				string job = params.get((Json::Value::ArrayIndex)0, "").asString();
-				string sHeaderHash = params.get((Json::Value::ArrayIndex)1, "").asString();
-				string sSeedHash = params.get((Json::Value::ArrayIndex)2, "").asString();
-				string sShareTarget = params.get((Json::Value::ArrayIndex)3, "").asString();
-				//bool cleanJobs = params.get((Json::Value::ArrayIndex)4, "").asBool();
-				
-				// coinmine.pl fix
-				int l = sShareTarget.length();
-				if (l < 66)
-					sShareTarget = "0x" + string(66 - l, '0') + sShareTarget.substr(2);
-								
-
-				if (sHeaderHash != "" && sSeedHash != "" && sShareTarget != "")
-				{
-					cnote << "Received new job #" + job.substr(0,8);
-					//cnote << "Header hash: " + sHeaderHash;
-					//cnote << "Seed hash: " + sSeedHash;
-					//cnote << "Share target: " + sShareTarget;
 
-					h256 seedHash = h256(sSeedHash);
-					h256 headerHash = h256(sHeaderHash);
+				if (m_ethereumStratum)
+				{
+					string sSeedHash = params.get((Json::Value::ArrayIndex)1, "").asString();
+					string sHeaderHash = params.get((Json::Value::ArrayIndex)2, "").asString();
 
-					if (headerHash != m_current.headerHash)
+					if (sHeaderHash != "" && sSeedHash != "")
 					{
-						//x_current.lock();
-						if (p_worktimer)
-							p_worktimer->cancel();
+						cnote << "Received new job #" + job;
+						//cnote << "Header hash: " + sHeaderHash;
+						//cnote << "Seed hash: " + sSeedHash;
+						//cnote << "Share target: " + sShareTarget;
+
+						h256 seedHash = h256(sSeedHash);
+						h256 headerHash = h256(sHeaderHash);
 
 						m_previous.headerHash = m_current.headerHash;
 						m_previous.seedHash = m_current.seedHash;
 						m_previous.boundary = m_current.boundary;
+						m_previous.startNonce = m_current.startNonce;
+						m_previous.exSizeBits = m_previous.exSizeBits;
 						m_previousJob = m_job;
 
 						m_current.headerHash = h256(sHeaderHash);
 						m_current.seedHash = seedHash;
-						m_current.boundary = h256(sShareTarget);// , h256::AlignRight);
+						m_current.boundary = h256();
+						diffToTarget((uint32_t*)m_current.boundary.data(), m_nextWorkDifficulty);
+						m_current.startNonce = ethash_swap_u64(*((uint64_t*)m_extraNonce.data()));
+						m_current.exSizeBits = m_extraNonceHexSize * 4;
 						m_job = job;
 
 						p_farm->setWork(m_current);
-						//x_current.unlock();
-						p_worktimer = new boost::asio::deadline_timer(m_io_service, boost::posix_time::seconds(m_worktimeout));
-						p_worktimer->async_wait(boost::bind(&EthStratumClient::work_timeout_handler, this, boost::asio::placeholders::error));
+					}
+				}
+				else
+				{
+					string sHeaderHash = params.get((Json::Value::ArrayIndex)1, "").asString();
+					string sSeedHash = params.get((Json::Value::ArrayIndex)2, "").asString();
+					string sShareTarget = params.get((Json::Value::ArrayIndex)3, "").asString();
+					//bool cleanJobs = params.get((Json::Value::ArrayIndex)4, "").asBool();
 
+					// coinmine.pl fix
+					int l = sShareTarget.length();
+					if (l < 66)
+						sShareTarget = "0x" + string(66 - l, '0') + sShareTarget.substr(2);
+
+
+					if (sHeaderHash != "" && sSeedHash != "" && sShareTarget != "")
+					{
+						cnote << "Received new job #" + job.substr(0, 8);
+						//cnote << "Header hash: " + sHeaderHash;
+						//cnote << "Seed hash: " + sSeedHash;
+						//cnote << "Share target: " + sShareTarget;
+
+						h256 seedHash = h256(sSeedHash);
+						h256 headerHash = h256(sHeaderHash);
+
+						if (headerHash != m_current.headerHash)
+						{
+							//x_current.lock();
+							if (p_worktimer)
+								p_worktimer->cancel();
+
+							m_previous.headerHash = m_current.headerHash;
+							m_previous.seedHash = m_current.seedHash;
+							m_previous.boundary = m_current.boundary;
+							m_previousJob = m_job;
+
+							m_current.headerHash = h256(sHeaderHash);
+							m_current.seedHash = seedHash;
+							m_current.boundary = h256(sShareTarget);// , h256::AlignRight);
+							m_job = job;
+
+							p_farm->setWork(m_current);
+							//x_current.unlock();
+							p_worktimer = new boost::asio::deadline_timer(m_io_service, boost::posix_time::seconds(m_worktimeout));
+							p_worktimer->async_wait(boost::bind(&EthStratumClient::work_timeout_handler, this, boost::asio::placeholders::error));
+
+						}
 					}
 				}
 			}
 		}
-		else if (method == "mining.set_difficulty")
+		else if (method == "mining.set_difficulty" && m_ethereumStratum)
 		{
-
+			params = responseObject.get("params", Json::Value::null);
+			if (params.isArray())
+			{
+				m_nextWorkDifficulty = params.get((Json::Value::ArrayIndex)0, 1).asDouble();
+				if (m_nextWorkDifficulty <= 0.0001) m_nextWorkDifficulty = 0.0001;
+				cnote << "Difficulty set to " << m_nextWorkDifficulty;
+			}
+		}
+		else if (method == "mining.set_extranonce" && m_ethereumStratum)
+		{
+			params = responseObject.get("params", Json::Value::null);
+			if (params.isArray())
+			{
+				std::string enonce = params.get((Json::Value::ArrayIndex)0, "").asString();
+				processExtranonce(enonce);
+			}
 		}
 		else if (method == "client.get_version")
 		{
@@ -364,11 +470,17 @@ bool EthStratumClient::submit(EthashProofOfWork::Solution solution) {
 	x_current.unlock();
 
 	cnote << "Solution found; Submitting to" << p_active->host << "...";
-	cnote << "  Nonce:" << "0x" + solution.nonce.hex();
+	string minernonce = solution.nonce.hex().substr(m_extraNonceHexSize, 16 - m_extraNonceHexSize);
+	if (!m_ethereumStratum)
+		cnote << "  Nonce:" << "0x" + solution.nonce.hex();
 
 	if (EthashAux::eval(tempWork.seedHash, tempWork.headerHash, solution.nonce).value < tempWork.boundary)
 	{
-		string json = "{\"id\": 4, \"method\": \"mining.submit\", \"params\": [\"" + p_active->user + "\",\"" + temp_job + "\",\"0x" + solution.nonce.hex() + "\",\"0x" + tempWork.headerHash.hex() + "\",\"0x" + solution.mixHash.hex() + "\"]}\n";
+		string json;
+		if (m_ethereumStratum)
+			json = "{\"id\": 4, \"method\": \"mining.submit\", \"params\": [\"" + p_active->user + "\",\"" + temp_job + "\",\"" + minernonce + "\"]}\n";
+		else
+			json = "{\"id\": 4, \"method\": \"mining.submit\", \"params\": [\"" + p_active->user + "\",\"" + temp_job + "\",\"0x" + solution.nonce.hex() + "\",\"0x" + tempWork.headerHash.hex() + "\",\"0x" + solution.mixHash.hex() + "\"]}\n";
 		std::ostream os(&m_requestBuffer);
 		os << json;
 		m_stale = false;
@@ -379,7 +491,11 @@ bool EthStratumClient::submit(EthashProofOfWork::Solution solution) {
 	}
 	else if (EthashAux::eval(tempPreviousWork.seedHash, tempPreviousWork.headerHash, solution.nonce).value < tempPreviousWork.boundary)
 	{
-		string json = "{\"id\": 4, \"method\": \"mining.submit\", \"params\": [\"" + p_active->user + "\",\"" + temp_previous_job + "\",\"0x" + solution.nonce.hex() + "\",\"0x" + tempPreviousWork.headerHash.hex() + "\",\"0x" + solution.mixHash.hex() + "\"]}\n";
+		string json;
+		if (m_ethereumStratum)
+			json = "{\"id\": 4, \"method\": \"mining.submit\", \"params\": [\"" + p_active->user + "\",\"" + temp_previous_job + "\",\"" + minernonce + "\"]}\n";
+		else
+			json = "{\"id\": 4, \"method\": \"mining.submit\", \"params\": [\"" + p_active->user + "\",\"" + temp_previous_job + "\",\"0x" + solution.nonce.hex() + "\",\"0x" + tempPreviousWork.headerHash.hex() + "\",\"0x" + solution.mixHash.hex() + "\"]}\n";
 		std::ostream os(&m_requestBuffer);
 		os << json;
 		m_stale = true;

libstratum/EthStratumClient.h

@@ -21,7 +21,7 @@ using namespace dev::eth;
 class EthStratumClient
 {
 public:
-	EthStratumClient(GenericFarm<EthashProofOfWork> * f, MinerType m, string const & host, string const & port, string const & user, string const & pass, int const & retries, int const & worktimeout);
+	EthStratumClient(GenericFarm<EthashProofOfWork> * f, MinerType m, string const & host, string const & port, string const & user, string const & pass, int const & retries, int const & worktimeout, bool ethstratum);
 	~EthStratumClient();
 
 	void setFailover(string const & host, string const & port);
@@ -87,4 +87,12 @@ private:
 
 	boost::asio::deadline_timer * p_worktimer;
 
+	bool m_ethereumStratum;
+
+	double m_nextWorkDifficulty;
+
+	h64 m_extraNonce;
+	int m_extraNonceHexSize;
+
+	void processExtranonce(std::string& enonce);
 };

libstratum/EthStratumClientV2.cpp

@@ -1,10 +1,33 @@
 
 #include "EthStratumClientV2.h"
 #include <libdevcore/Log.h>
+#include <libethash/endian.h>
 using boost::asio::ip::tcp;
 
 
-EthStratumClientV2::EthStratumClientV2(GenericFarm<EthashProofOfWork> * f, MinerType m, string const & host, string const & port, string const & user, string const & pass, int const & retries, int const & worktimeout)
+static void diffToTarget(uint32_t *target, double diff)
+{
+	uint32_t target2[8];
+	uint64_t m;
+	int k;
+
+	for (k = 6; k > 0 && diff > 1.0; k--)
+		diff /= 4294967296.0;
+	m = (uint64_t)(4294901760.0 / diff);
+	if (m == 0 && k == 6)
+		memset(target2, 0xff, 32);
+	else {
+		memset(target2, 0, 32);
+		target2[k] = (uint32_t)m;
+		target2[k + 1] = (uint32_t)(m >> 32);
+	}
+
+	for (int i = 0; i < 32; i++)
+		((uint8_t*)target)[31 - i] = ((uint8_t*)target2)[i];
+}
+
+
+EthStratumClientV2::EthStratumClientV2(GenericFarm<EthashProofOfWork> * f, MinerType m, string const & host, string const & port, string const & user, string const & pass, int const & retries, int const & worktimeout, bool ethstratum)
 	: Worker("stratum"), 
 	  m_socket(m_io_service)
 {
@@ -21,6 +44,8 @@ EthStratumClientV2::EthStratumClientV2(GenericFarm<EthashProofOfWork> * f, Miner
 	m_maxRetries = retries;
 	m_worktimeout = worktimeout;
 
+	m_ethereumStratum = ethstratum;
+
 	p_farm = f;
 	p_worktimer = nullptr;
 	startWorking();
@@ -127,7 +152,10 @@ void EthStratumClientV2::connect()
 			}
 		}
 		std::ostream os(&m_requestBuffer);
-		os << "{\"id\": 1, \"method\": \"mining.subscribe\", \"params\": []}\n";
+		if (m_ethereumStratum)
+			os << "{\"id\": 1, \"method\": \"mining.subscribe\", \"params\": [\"ethminer/" << ETH_PROJECT_VERSION << "\",\"EthereumStratum/1.0.0\"]}\n";
+		else
+			os << "{\"id\": 1, \"method\": \"mining.subscribe\", \"params\": []}\n";
 		write(m_socket, m_requestBuffer);
 	}
 }
@@ -184,6 +212,16 @@ void EthStratumClientV2::disconnect()
 	//m_io_service.stop();
 }
 
+void EthStratumClientV2::processExtranonce(std::string& enonce)
+{
+	m_extraNonceHexSize = enonce.length();
+
+	cnote << "Extranonce set to " << enonce;
+
+	for (int i = enonce.length(); i < 16; ++i) enonce += "0";
+	m_extraNonce = h64(enonce);
+}
+
 void EthStratumClientV2::processReponse(Json::Value& responseObject)
 {
 	Json::Value error = responseObject.get("error", new Json::Value);
@@ -200,11 +238,27 @@ void EthStratumClientV2::processReponse(Json::Value& responseObject)
 	case 1:
 		cnote << "Subscribed to stratum server";
 
-		os << "{\"id\": 2, \"method\": \"mining.authorize\", \"params\": [\"" << p_active->user << "\",\"" << p_active->pass << "\"]}\n";
+		if (m_ethereumStratum)
+		{
+			m_nextWorkDifficulty = 1;
+			params = responseObject.get("result", Json::Value::null);
+			if (params.isArray())
+			{
+				std::string enonce = params.get((Json::Value::ArrayIndex)1, "").asString();
+				processExtranonce(enonce);
+			}
+
+			os << "{\"id\": 2, \"method\": \"mining.extranonce.subscribe\", \"params\": []}\n";
+		}
+
+		os << "{\"id\": 3, \"method\": \"mining.authorize\", \"params\": [\"" << p_active->user << "\",\"" << p_active->pass << "\"]}\n";
 
 		write(m_socket, m_requestBuffer);
 		break;
 	case 2:
+		// nothing to do...
+		break;
+	case 3:
 		m_authorized = responseObject.get("result", Json::Value::null).asBool();
 		if (!m_authorized)
 		{
@@ -232,51 +286,104 @@ void EthStratumClientV2::processReponse(Json::Value& responseObject)
 			if (params.isArray())
 			{
 				string job = params.get((Json::Value::ArrayIndex)0, "").asString();
-				string sHeaderHash = params.get((Json::Value::ArrayIndex)1, "").asString();
-				string sSeedHash = params.get((Json::Value::ArrayIndex)2, "").asString();
-				string sShareTarget = params.get((Json::Value::ArrayIndex)3, "").asString();
-				//bool cleanJobs = params.get((Json::Value::ArrayIndex)4, "").asBool();
-				
-				// coinmine.pl fix
-				int l = sShareTarget.length();
-				if (l < 66)
-					sShareTarget = "0x" + string(66 - l, '0') + sShareTarget.substr(2);
-								
 
-				if (sHeaderHash != "" && sSeedHash != "" && sShareTarget != "")
+				if (m_ethereumStratum)
 				{
-					cnote << "Received new job #" + job.substr(0,8);
+					string job = params.get((Json::Value::ArrayIndex)0, "").asString();
+					string sSeedHash = params.get((Json::Value::ArrayIndex)1, "").asString();
+					string sHeaderHash = params.get((Json::Value::ArrayIndex)2, "").asString();
 
-					h256 seedHash = h256(sSeedHash);
-					h256 headerHash = h256(sHeaderHash);
-
-					if (headerHash != m_current.headerHash)
+					if (sHeaderHash != "" && sSeedHash != "")
 					{
-						//x_current.lock();
-						//if (p_worktimer)
-						//	p_worktimer->cancel();
+						cnote << "Received new job #" + job;
+						//cnote << "Header hash: " + sHeaderHash;
+						//cnote << "Seed hash: " + sSeedHash;
+						//cnote << "Share target: " + sShareTarget;
+
+						h256 seedHash = h256(sSeedHash);
+						h256 headerHash = h256(sHeaderHash);
 
 						m_previous.headerHash = m_current.headerHash;
 						m_previous.seedHash = m_current.seedHash;
 						m_previous.boundary = m_current.boundary;
+						m_previous.startNonce = m_current.startNonce;
+						m_previous.exSizeBits = m_previous.exSizeBits;
 						m_previousJob = m_job;
 
 						m_current.headerHash = h256(sHeaderHash);
 						m_current.seedHash = seedHash;
-						m_current.boundary = h256(sShareTarget);
+						m_current.boundary = h256();
+						diffToTarget((uint32_t*)m_current.boundary.data(), m_nextWorkDifficulty);
+						m_current.startNonce = ethash_swap_u64(*((uint64_t*)m_extraNonce.data()));
+						m_current.exSizeBits = m_extraNonceHexSize * 4;
 						m_job = job;
 
 						p_farm->setWork(m_current);
-						//x_current.unlock();
-						//p_worktimer = new boost::asio::deadline_timer(m_io_service, boost::posix_time::seconds(m_worktimeout));
-						//p_worktimer->async_wait(boost::bind(&EthStratumClientV2::work_timeout_handler, this, boost::asio::placeholders::error));
+					}
+				}
+				else
+				{
+					string sHeaderHash = params.get((Json::Value::ArrayIndex)1, "").asString();
+					string sSeedHash = params.get((Json::Value::ArrayIndex)2, "").asString();
+					string sShareTarget = params.get((Json::Value::ArrayIndex)3, "").asString();
+					//bool cleanJobs = params.get((Json::Value::ArrayIndex)4, "").asBool();
+
+					// coinmine.pl fix
+					int l = sShareTarget.length();
+					if (l < 66)
+						sShareTarget = "0x" + string(66 - l, '0') + sShareTarget.substr(2);
+
+
+					if (sHeaderHash != "" && sSeedHash != "" && sShareTarget != "")
+					{
+						cnote << "Received new job #" + job.substr(0, 8);
+
+						h256 seedHash = h256(sSeedHash);
+						h256 headerHash = h256(sHeaderHash);
+
+						if (headerHash != m_current.headerHash)
+						{
+							//x_current.lock();
+							//if (p_worktimer)
+							//	p_worktimer->cancel();
+
+							m_previous.headerHash = m_current.headerHash;
+							m_previous.seedHash = m_current.seedHash;
+							m_previous.boundary = m_current.boundary;
+							m_previousJob = m_job;
+
+							m_current.headerHash = h256(sHeaderHash);
+							m_current.seedHash = seedHash;
+							m_current.boundary = h256(sShareTarget);
+							m_job = job;
+
+							p_farm->setWork(m_current);
+							//x_current.unlock();
+							//p_worktimer = new boost::asio::deadline_timer(m_io_service, boost::posix_time::seconds(m_worktimeout));
+							//p_worktimer->async_wait(boost::bind(&EthStratumClientV2::work_timeout_handler, this, boost::asio::placeholders::error));
+						}
 					}
 				}
 			}
 		}
-		else if (method == "mining.set_difficulty")
+		else if (method == "mining.set_difficulty" && m_ethereumStratum)
 		{
-
+			params = responseObject.get("params", Json::Value::null);
+			if (params.isArray())
+			{
+				m_nextWorkDifficulty = params.get((Json::Value::ArrayIndex)0, 1).asDouble();
+				if (m_nextWorkDifficulty <= 0.0001) m_nextWorkDifficulty = 0.0001;
+				cnote << "Difficulty set to " << m_nextWorkDifficulty;
+			}
+		}
+		else if (method == "mining.set_extranonce" && m_ethereumStratum)
+		{
+			params = responseObject.get("params", Json::Value::null);
+			if (params.isArray())
+			{
+				std::string enonce = params.get((Json::Value::ArrayIndex)0, "").asString();
+				processExtranonce(enonce);
+			}
 		}
 		else if (method == "client.get_version")
 		{
@@ -304,11 +411,17 @@ bool EthStratumClientV2::submit(EthashProofOfWork::Solution solution) {
 	x_current.unlock();
 
 	cnote << "Solution found; Submitting to" << p_active->host << "...";
-	cnote << "  Nonce:" << "0x" + solution.nonce.hex();
+	string minernonce = solution.nonce.hex().substr(m_extraNonceHexSize, 16 - m_extraNonceHexSize);
+	if (!m_ethereumStratum)
+		cnote << "  Nonce:" << "0x" + solution.nonce.hex();
 
 	if (EthashAux::eval(tempWork.seedHash, tempWork.headerHash, solution.nonce).value < tempWork.boundary)
 	{
-		string json = "{\"id\": 4, \"method\": \"mining.submit\", \"params\": [\"" + p_active->user + "\",\"" + temp_job + "\",\"0x" + solution.nonce.hex() + "\",\"0x" + tempWork.headerHash.hex() + "\",\"0x" + solution.mixHash.hex() + "\"]}\n";
+		string json;
+		if (m_ethereumStratum)
+			json = "{\"id\": 4, \"method\": \"mining.submit\", \"params\": [\"" + p_active->user + "\",\"" + temp_job + "\",\"" + minernonce + "\"]}\n";
+		else
+			json = "{\"id\": 4, \"method\": \"mining.submit\", \"params\": [\"" + p_active->user + "\",\"" + temp_job + "\",\"0x" + solution.nonce.hex() + "\",\"0x" + tempWork.headerHash.hex() + "\",\"0x" + solution.mixHash.hex() + "\"]}\n";
 		std::ostream os(&m_requestBuffer);
 		os << json;
 		m_stale = false;
@@ -317,7 +430,11 @@ bool EthStratumClientV2::submit(EthashProofOfWork::Solution solution) {
 	}
 	else if (EthashAux::eval(tempPreviousWork.seedHash, tempPreviousWork.headerHash, solution.nonce).value < tempPreviousWork.boundary)
 	{
-		string json = "{\"id\": 4, \"method\": \"mining.submit\", \"params\": [\"" + p_active->user + "\",\"" + temp_previous_job + "\",\"0x" + solution.nonce.hex() + "\",\"0x" + tempPreviousWork.headerHash.hex() + "\",\"0x" + solution.mixHash.hex() + "\"]}\n";
+		string json;
+		if (m_ethereumStratum)
+			json = "{\"id\": 4, \"method\": \"mining.submit\", \"params\": [\"" + p_active->user + "\",\"" + temp_previous_job + "\",\"" + minernonce + "\"]}\n";
+		else
+			json = "{\"id\": 4, \"method\": \"mining.submit\", \"params\": [\"" + p_active->user + "\",\"" + temp_previous_job + "\",\"0x" + solution.nonce.hex() + "\",\"0x" + tempPreviousWork.headerHash.hex() + "\",\"0x" + solution.mixHash.hex() + "\"]}\n";
 		std::ostream os(&m_requestBuffer);
 		os << json;
 		m_stale = true;

libstratum/EthStratumClientV2.h

@@ -22,7 +22,7 @@ using namespace dev::eth;
 class EthStratumClientV2 : public Worker
 {
 public:
-	EthStratumClientV2(GenericFarm<EthashProofOfWork> * f, MinerType m, string const & host, string const & port, string const & user, string const & pass, int const & retries, int const & worktimeout);
+	EthStratumClientV2(GenericFarm<EthashProofOfWork> * f, MinerType m, string const & host, string const & port, string const & user, string const & pass, int const & retries, int const & worktimeout, bool ethstratum);
 	~EthStratumClientV2();
 
 	void setFailover(string const & host, string const & port);
@@ -81,4 +81,12 @@ private:
 
 	boost::asio::deadline_timer * p_worktimer;
 
+	bool m_ethereumStratum;
+
+	double m_nextWorkDifficulty;
+
+	h64 m_extraNonce;
+	int m_extraNonceHexSize;
+
+	void processExtranonce(std::string& enonce);
 };