Merge branch '108'

Conflicts:
	libethash-cl/ethash_cl_miner.cpp

CMakeLists.txt

@@ -2,7 +2,7 @@
 cmake_minimum_required(VERSION 2.8.12)
 
 set(PROJECT_VERSION "0.9.41")
-set(GENOIL_VERSION "1.0.7")
+set(GENOIL_VERSION "1.0.8")
 if (${CMAKE_VERSION} VERSION_GREATER 3.0)
 	cmake_policy(SET CMP0042 OLD) 	# fix MACOSX_RPATH
 	cmake_policy(SET CMP0048 NEW) 	# allow VERSION argument in project()

ethminer/MinerAux.h

@@ -90,7 +90,7 @@ inline std::string credits()
 class BadArgument: public Exception {};
 struct MiningChannel: public LogChannel
 {
-	static const char* name() { return EthGreen "miner"; }
+	static const char* name() { return EthGreen "  m"; }
 	static const int verbosity = 2;
 	static const bool debug = false;
 };
@@ -145,6 +145,7 @@ public:
 		}
 		else if (arg == "--farm-recheck" && i + 1 < argc)
 			try {
+				m_farmRecheckSet = true;
 				m_farmRecheckPeriod = stol(argv[++i]);
 			}
 			catch (...)
@@ -218,6 +219,11 @@ public:
 		{
 			m_fport = string(argv[++i]);
 		}
+		else if ((arg == "--work-timeout") && i + 1 < argc)
+		{
+			m_worktimeout = atoi(argv[++i]);
+		}
+		
 #endif
 		else if (arg == "--opencl-platform" && i + 1 < argc)
 			try {
@@ -228,17 +234,7 @@ public:
 				cerr << "Bad " << arg << " option: " << argv[i] << endl;
 				BOOST_THROW_EXCEPTION(BadArgument());
 			}
-		else if (arg == "--opencl-device" && i + 1 < argc)
-			try {
-				m_openclDevice = stol(argv[++i]);
-				m_miningThreads = 1;
-			}
-			catch (...)
-			{
-				cerr << "Bad " << arg << " option: " << argv[i] << endl;
-				BOOST_THROW_EXCEPTION(BadArgument());
-			}
-		else if (arg == "--opencl-devices")
+		else if (arg == "--opencl-devices" || arg == "--opencl-device")
 			while (m_openclDeviceCount < 16 && i + 1 < argc)
 			{
 				try
@@ -598,6 +594,7 @@ public:
 			<< "	-FS, --failover-stratum <host:port>  Failover stratum server at host:port" << endl
 			<< "    -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 60. Don't set lower than max. avg. block time" << 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
@@ -875,7 +872,8 @@ private:
 			f.start("opencl");
 		else if (_m == MinerType::CUDA)
 			f.start("cuda");
-		EthashProofOfWork::WorkPackage current;
+		EthashProofOfWork::WorkPackage current, previous;
+		boost::mutex x_current;
 		EthashAux::FullType dag;
 		while (m_running)
 			try
@@ -892,7 +890,7 @@ private:
 					auto mp = f.miningProgress();
 					f.resetMiningProgress();
 					if (current)
-						minelog << "Mining on PoWhash" << current.headerHash << ": " << mp;
+						minelog << "Mining on PoWhash" << "#" + (current.headerHash.hex().substr(0, 8)) << ": " << mp << f.getSolutionStats();
 					else
 						minelog << "Getting work package...";
 
@@ -912,41 +910,69 @@ private:
 					h256 hh(v[0].asString());
 					h256 newSeedHash(v[1].asString());
 					if (current.seedHash != newSeedHash)
+					{
 						minelog << "Grabbing DAG for" << newSeedHash;
+					}
 					if (!(dag = EthashAux::full(newSeedHash, true, [&](unsigned _pc){ cout << "\rCreating DAG. " << _pc << "% done..." << flush; return 0; })))
+					{
 						BOOST_THROW_EXCEPTION(DAGCreationFailure());
+					}
 					if (m_precompute)
+					{
 						EthashAux::computeFull(sha3(newSeedHash), true);
+					}
 					if (hh != current.headerHash)
 					{
+						x_current.lock();
+						previous.headerHash = current.headerHash;
+						previous.seedHash = current.seedHash;
+						previous.boundary = current.boundary;
 						current.headerHash = hh;
 						current.seedHash = newSeedHash;
 						current.boundary = h256(fromHex(v[2].asString()), h256::AlignRight);
-						minelog << "Got work package:";
-						minelog << "  Header-hash:" << current.headerHash.hex();
-						minelog << "  Seedhash:" << current.seedHash.hex();
-						minelog << "  Target: " << h256(current.boundary).hex();
+						minelog << "Got work package: #" + current.headerHash.hex().substr(0,8);
+						//minelog << "  Seedhash:" << current.seedHash.hex();
+						//minelog << "  Target: " << h256(current.boundary).hex();
 						f.setWork(current);
+						x_current.unlock();
 					}
 					this_thread::sleep_for(chrono::milliseconds(_recheckPeriod));
 				}
 				cnote << "Solution found; Submitting to" << _remote << "...";
 				cnote << "  Nonce:" << solution.nonce.hex();
-				cnote << "  Mixhash:" << solution.mixHash.hex();
-				cnote << "  Header-hash:" << current.headerHash.hex();
-				cnote << "  Seedhash:" << current.seedHash.hex();
-				cnote << "  Target: " << h256(current.boundary).hex();
-				cnote << "  Ethash: " << h256(EthashAux::eval(current.seedHash, current.headerHash, solution.nonce).value).hex();
+				//cnote << "  Mixhash:" << solution.mixHash.hex();
+				//cnote << "  Header-hash:" << current.headerHash.hex();
+				//cnote << "  Seedhash:" << solved.seedHash.hex();
+				//cnote << "  Target: " << h256(solved.boundary).hex();
+				//cnote << "  Ethash: " << h256(EthashAux::eval(solved.seedHash, solved.headerHash, solution.nonce).value).hex();
 				if (EthashAux::eval(current.seedHash, current.headerHash, solution.nonce).value < current.boundary)
 				{
 					bool ok = prpc->eth_submitWork("0x" + toString(solution.nonce), "0x" + toString(current.headerHash), "0x" + toString(solution.mixHash));
-					if (ok)
+					if (ok) {
 						cnote << "B-) Submitted and accepted.";
-					else
+						f.acceptedSolution(false);
+					}
+					else {
 						cwarn << ":-( Not accepted.";
+						f.rejectedSolution(false);
+					}
 				}
-				else
+				else if (EthashAux::eval(previous.seedHash, previous.headerHash, solution.nonce).value < previous.boundary)
+				{
+					bool ok = prpc->eth_submitWork("0x" + toString(solution.nonce), "0x" + toString(previous.headerHash), "0x" + toString(solution.mixHash));
+					if (ok) {
+						cnote << "B-) Submitted and accepted.";
+						f.acceptedSolution(true);
+					}
+					else {
+						cwarn << ":-( Not accepted.";
+						f.rejectedSolution(true);
+					}
+				}
+				else {
+					f.failedSolution();
 					cwarn << "FAILURE: GPU gave incorrect result!";
+				}
 				current.reset();
 			}
 			catch (jsonrpc::JsonRpcException&)
@@ -998,8 +1024,11 @@ private:
 #if ETH_ETHASHCUDA
 		sealers["cuda"] = GenericFarm<EthashProofOfWork>::SealerDescriptor{ &EthashCUDAMiner::instances, [](GenericMiner<EthashProofOfWork>::ConstructionInfo ci){ return new EthashCUDAMiner(ci); } };
 #endif
+		if (!m_farmRecheckSet)
+			m_farmRecheckPeriod = m_defaultStratumFarmRecheckPeriod;
+		
 		GenericFarm<EthashProofOfWork> f;
-		EthStratumClient client(&f, m_minerType, m_farmURL, m_port, m_user, m_pass, m_maxFarmRetries, m_precompute);
+		EthStratumClient client(&f, m_minerType, m_farmURL, m_port, m_user, m_pass, m_maxFarmRetries, m_worktimeout, m_precompute);
 		if (m_farmFailOverURL != "")
 		{
 			if (m_fuser != "")
@@ -1026,8 +1055,8 @@ private:
 			if (client.isConnected())
 			{
 				if (client.current())
-					minelog << "Mining on PoWhash" << client.currentHeaderHash() << ": " << mp;
-				else
+					minelog << "Mining on PoWhash" << "#"+(client.currentHeaderHash().hex().substr(0,8)) << ": " << mp << f.getSolutionStats();
+				else if (client.waitState() == MINER_WAIT_STATE_WORK)
 					minelog << "Waiting for work package...";
 			}
 			this_thread::sleep_for(chrono::milliseconds(m_farmRecheckPeriod));
@@ -1086,6 +1115,9 @@ private:
 	unsigned m_farmRetries = 0;
 	unsigned m_maxFarmRetries = 3;
 	unsigned m_farmRecheckPeriod = 500;
+	unsigned m_defaultStratumFarmRecheckPeriod = 2000;
+	bool m_farmRecheckSet = false;
+	int m_worktimeout = 90;
 	bool m_precompute = true;
 
 #if ETH_STRATUM || !ETH_TRUE

ethminer/main.cpp

@@ -72,7 +72,7 @@ int main(int argc, char** argv)
 	cout << "=====================================================================" << endl;
 	cout << "Forked from github.com/ethereum/cpp-ethereum" << endl;
 	cout << "CUDA kernel ported from Tim Hughes' OpenCL kernel" << endl;
-	cout << "With contributions from RoBiK, tpruvot and sp_ " << endl << endl;
+	cout << "With contributions from nerdralph, RoBiK, tpruvot and sp_ " << endl << endl;
 	cout << "Please consider a donation to:" << endl;
 	cout << "ETH: 0xeb9310b185455f863f526dab3d245809f6854b4d" << endl << endl;
 

libdevcore/Common.cpp

@@ -54,7 +54,9 @@ TimerHelper::~TimerHelper()
 {
 	auto e = std::chrono::high_resolution_clock::now() - m_t;
 	if (!m_ms || e > chrono::milliseconds(m_ms))
+	{
 		clog(TimerChannel) << m_id << chrono::duration_cast<chrono::milliseconds>(e).count() << "ms";
+	}
 }
 
 }

libdevcore/TransientDirectory.cpp

@@ -60,5 +60,7 @@ TransientDirectory::~TransientDirectory()
 	ec.clear();
 	fs::remove_all(m_path, ec);
 	if (!ec)
+	{
 		cwarn << "Failed to delete directory '" << m_path << "': " << ec.message();
+	}
 }

libethash-cl/ethash_cl_miner.cpp

@@ -52,6 +52,15 @@
 #define CL_MEM_HOST_READ_ONLY 0
 #endif
 
+// apple fix
+#ifndef CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV
+#define CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV       0x4000
+#endif
+
+#ifndef CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV
+#define CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV       0x4001
+#endif
+
 #undef min
 #undef max
 
@@ -379,7 +388,7 @@ bool ethash_cl_miner::init(
 			sprintf(options, "-cl-nv-maxrregcount=%d", maxregs);// , computeCapability);
 		}
 		else {
-			sprintf(options, "");
+			sprintf(options, "%s", "");
 		}
 		// create context
 		m_context = cl::Context(vector<cl::Device>(&device, &device + 1));
@@ -401,10 +410,6 @@ bool ethash_cl_miner::init(
 		addDefinition(code, "PLATFORM", platformId);
 		addDefinition(code, "COMPUTE", computeCapability);
 
-		//debugf("%s", code.c_str());
-
-
-		
 		// create miner OpenCL program
 		cl::Program::Sources sources;
 		sources.push_back({ code.c_str(), code.size() });
@@ -441,6 +446,10 @@ bool ethash_cl_miner::init(
 		ETHCL_LOG("Creating buffer for header.");
 		m_header = cl::Buffer(m_context, CL_MEM_READ_ONLY, 32);
 
+		m_searchKernel.setArg(1, m_header);
+		m_searchKernel.setArg(2, m_dag);
+		m_searchKernel.setArg(5, ~0u);
+
 		// create mining buffers
 		for (unsigned i = 0; i != c_bufferCount; ++i)
 		{
@@ -456,15 +465,16 @@ bool ethash_cl_miner::init(
 	return true;
 }
 
+typedef struct 
+{
+	uint64_t start_nonce;
+	unsigned buf;
+} pending_batch;
+
 void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook& hook)
 {
 	try
 	{
-		struct pending_batch
-		{
-			uint64_t start_nonce;
-			unsigned buf;
-		};
 		queue<pending_batch> pending;
 
 		// this can't be a static because in MacOSX OpenCL implementation a segfault occurs when a static is passed to OpenCL functions
@@ -483,19 +493,14 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook
 #endif
 			m_queue.finish();
 
-		
-		m_searchKernel.setArg(1, m_header);
-		m_searchKernel.setArg(2, m_dag );
 		// pass these to stop the compiler unrolling the loops
 		m_searchKernel.setArg(4, target);
-		m_searchKernel.setArg(5, ~0u);
-
+		
 		unsigned buf = 0;
 		random_device engine;
 		uint64_t start_nonce = uniform_int_distribution<uint64_t>()(engine);
 		for (;; start_nonce += m_globalWorkSize)
 		{
-			auto t = chrono::high_resolution_clock::now();
 			// supply output buffer to kernel
 			m_searchKernel.setArg(0, m_searchBuffer[buf]);
 			m_searchKernel.setArg(3, start_nonce);

libethash-cl/ethash_cl_miner.h

@@ -6,10 +6,10 @@
 #if defined(__clang__)
 #pragma clang diagnostic push
 #pragma clang diagnostic ignored "-Wunused-parameter"
-#include "cl.hpp"
+#include "CL/cl.hpp"
 #pragma clang diagnostic pop
 #else
-#include "cl.hpp"
+#include "CL/cl.hpp"
 #endif
 
 #include <time.h>

libethash-cuda/CMakeLists.txt

@@ -13,7 +13,7 @@ LIST(APPEND CUDA_NVCC_FLAGS_DEBUG -G)
 if(COMPUTE AND (COMPUTE GREATER 0))
 	LIST(APPEND CUDA_NVCC_FLAGS -gencode arch=compute_${COMPUTE},code=sm_${COMPUTE})
 else(COMPUTE AND (COMPUTE GREATER 0))
-	set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-gencode arch=compute_30,code=sm_30;-gencode arch=compute_35,code=sm_35;-gencode arch=compute_50,code=sm_50;-gencode arch=compute_52,code=sm_52)
+	set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-gencode arch=compute_50,code=sm_50;-gencode arch=compute_52,code=sm_52)
 endif(COMPUTE AND (COMPUTE GREATER 0))
 
 

libethash-cuda/ethash_cuda_miner_kernel.cu

@@ -8,7 +8,96 @@
 #include "ethash_cuda_miner_kernel_globals.h"
 #include "cuda_helper.h"
 
-#include "dagger_shuffled.cuh"
+#include "keccak.cuh"
+#include "fnv.cuh"
+
+#define ACCESSES 64
+#define THREADS_PER_HASH (128 / 16)
+
+__device__ uint64_t compute_hash_shuffle(
+	uint64_t nonce
+	)
+{
+	// sha3_512(header .. nonce)
+	uint2 state[25];
+
+	state[4] = vectorize(nonce);
+
+	keccak_f1600_init(state);
+
+	// Threads work together in this phase in groups of 8.
+	const int thread_id = threadIdx.x &  (THREADS_PER_HASH - 1);
+	const int start_lane = threadIdx.x & ~(THREADS_PER_HASH - 1);
+	const int mix_idx = thread_id & 3;
+
+	uint4 mix;
+	uint2 shuffle[8];
+
+	for (int i = 0; i < THREADS_PER_HASH; i++)
+	{
+		// share init among threads
+		for (int j = 0; j < 8; j++) {
+			shuffle[j].x = __shfl(state[j].x, i, THREADS_PER_HASH);
+			shuffle[j].y = __shfl(state[j].y, i, THREADS_PER_HASH);
+		}
+
+		// ugly but avoids local reads/writes
+		if (mix_idx < 2) {
+			if (mix_idx == 0)
+				mix = vectorize2(shuffle[0], shuffle[1]);
+			else
+				mix = vectorize2(shuffle[2], shuffle[3]);
+		}
+		else  {
+			if (mix_idx == 2)
+				mix = vectorize2(shuffle[4], shuffle[5]);
+			else
+				mix = vectorize2(shuffle[6], shuffle[7]);
+		}
+
+		uint32_t init0 = __shfl(shuffle[0].x, start_lane);
+
+		for (uint32_t a = 0; a < ACCESSES; a += 4)
+		{
+			int t = ((a >> 2) & (THREADS_PER_HASH - 1));
+
+			for (uint32_t b = 0; b < 4; b++)
+			{
+				if (thread_id == t)
+				{
+					shuffle[0].x = fnv(init0 ^ (a + b), ((uint32_t *)&mix)[b]) % d_dag_size;
+				}
+				shuffle[0].x = __shfl(shuffle[0].x, t, THREADS_PER_HASH);
+
+				mix = fnv4(mix, (&d_dag[shuffle[0].x])->uint4s[thread_id]);
+			}
+		}
+
+		uint32_t thread_mix = fnv_reduce(mix);
+
+		// update mix accross threads
+
+		shuffle[0].x = __shfl(thread_mix, 0, THREADS_PER_HASH);
+		shuffle[0].y = __shfl(thread_mix, 1, THREADS_PER_HASH);
+		shuffle[1].x = __shfl(thread_mix, 2, THREADS_PER_HASH);
+		shuffle[1].y = __shfl(thread_mix, 3, THREADS_PER_HASH);
+		shuffle[2].x = __shfl(thread_mix, 4, THREADS_PER_HASH);
+		shuffle[2].y = __shfl(thread_mix, 5, THREADS_PER_HASH);
+		shuffle[3].x = __shfl(thread_mix, 6, THREADS_PER_HASH);
+		shuffle[3].y = __shfl(thread_mix, 7, THREADS_PER_HASH);
+
+		if (i == thread_id) {
+			//move mix into state:
+			state[8] = shuffle[0];
+			state[9] = shuffle[1];
+			state[10] = shuffle[2];
+			state[11] = shuffle[3];
+		}
+	}
+
+	// keccak_256(keccak_512(header..nonce) .. mix);
+	return keccak_f1600_final(state);
+}
 
 __global__ void 
 __launch_bounds__(896, 1)

libethcore/EthashCUDAMiner.cpp

@@ -149,7 +149,8 @@ void EthashCUDAMiner::workLoop()
 	// take local copy of work since it may end up being overwritten by kickOff/pause.
 	try {
 		WorkPackage w = work();
-		cnote << "workLoop" << !!m_miner << m_minerSeed << w.seedHash;
+		//cnote << "seedhash" << "#" + m_minerSeed.hex().substr(0, 16);
+		cnote << "set work to" << "#" + w.headerHash.hex().substr(0, 8) + ", target " << "#" + w.boundary.hex().substr(0, 16);
 		if (!m_miner || m_minerSeed != w.seedHash)
 		{
 			cnote << "Initialising miner...";

libethcore/EthashGPUMiner.cpp

@@ -139,7 +139,7 @@ void EthashGPUMiner::workLoop()
 	// take local copy of work since it may end up being overwritten by kickOff/pause.
 	try {
 		WorkPackage w = work();
-		cnote << "workLoop" << !!m_miner << m_minerSeed << w.seedHash;
+		cnote << "set work to:" << "#" + w.headerHash.hex().substr(0, 8) + ", target " << "#" + w.boundary.hex().substr(0, 16);
 		if (!m_miner || m_minerSeed != w.seedHash)
 		{
 			cnote << "Initialising miner...";

libethcore/Farm.h

@@ -146,6 +146,36 @@ public:
 		resetTimer();
 	}
 
+	SolutionStats getSolutionStats() {
+		return m_solutionStats;
+	}
+
+	void failedSolution() {
+		m_solutionStats.failed();
+	}
+
+	void acceptedSolution(bool _stale) {
+		if (!_stale)
+		{
+			m_solutionStats.accepted();
+		}
+		else
+		{
+			m_solutionStats.acceptedStale();
+		}
+	}
+
+	void rejectedSolution(bool _stale) {
+		if (!_stale)
+		{
+			m_solutionStats.rejected();
+		}
+		else
+		{
+			m_solutionStats.rejectedStale();
+		}
+	}
+
 	using SolutionFound = std::function<bool(Solution const&)>;
 
 	/**
@@ -200,7 +230,11 @@ private:
 
 	std::map<std::string, SealerDescriptor> m_sealers;
 	std::string m_lastSealer;
-};
+
+	mutable SharedMutex x_solutionStats;
+	mutable SolutionStats m_solutionStats;
+
+}; 
 
 }
 }

libethcore/Miner.h

@@ -30,6 +30,10 @@
 #include <libdevcore/Worker.h>
 #include <libethcore/Common.h>
 
+#define MINER_WAIT_STATE_UNKNOWN 0
+#define MINER_WAIT_STATE_WORK	 1
+#define MINER_WAIT_STATE_DAG	 2
+
 namespace dev
 {
 
@@ -47,10 +51,44 @@ struct MineInfo: public WorkingProgress {};
 
 inline std::ostream& operator<<(std::ostream& _out, WorkingProgress _p)
 {
-	_out << _p.rate() << " H/s = " <<  _p.hashes << " hashes / " << (double(_p.ms) / 1000) << " s";
+	float mh = _p.rate() / 1000000.0f;
+	char mhs[16];
+	sprintf(mhs, "%.2f", mh);
+	_out << std::string(mhs) + "MH/s";
 	return _out;
 }
 
+class SolutionStats {
+public:
+	void accepted() { accepts++;  }
+	void rejected() { rejects++;  }
+	void failed()   { failures++; }
+
+	void acceptedStale() { acceptedStales++; }
+	void rejectedStale() { rejectedStales++; }
+
+
+	void reset() { accepts = rejects = failures = acceptedStales = rejectedStales = 0; }
+
+	unsigned getAccepts()			{ return accepts; }
+	unsigned getRejects()			{ return rejects; }
+	unsigned getFailures()			{ return failures; }
+	unsigned getAcceptedStales()	{ return acceptedStales; }
+	unsigned getRejectedStales()	{ return rejectedStales; }
+private:
+	unsigned accepts  = 0;
+	unsigned rejects  = 0;
+	unsigned failures = 0; 
+
+	unsigned acceptedStales = 0;
+	unsigned rejectedStales = 0;
+};
+
+inline std::ostream& operator<<(std::ostream& os, SolutionStats s)
+{
+	return os << "[A" << s.getAccepts() << "+" << s.getAcceptedStales() << ":R" << s.getRejects() << "+" << s.getRejectedStales() << ":F" << s.getFailures() << "]";
+}
+
 template <class PoW> class GenericMiner;
 
 /**

libstratum/EthStratumClient.cpp

@@ -1,10 +1,10 @@
 
 #include "EthStratumClient.h"
-
+#include <libdevcore/Log.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, bool const & precompute)
+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 const & precompute)
 	: m_socket(m_io_service)
 {
 	m_minerType = m;
@@ -20,8 +20,10 @@ EthStratumClient::EthStratumClient(GenericFarm<EthashProofOfWork> * f, MinerType
 	m_precompute = precompute;
 	m_pending = 0;
 	m_maxRetries = retries;
+	m_worktimeout = worktimeout;
 
 	p_farm = f;
+	p_worktimer = nullptr;
 	connect();
 }
 
@@ -70,11 +72,18 @@ void EthStratumClient::reconnect()
 		p_farm->stop();
 	}
 	*/
+	if (p_worktimer) {
+		p_worktimer->cancel();
+		p_worktimer = nullptr;
+	}
+
 	m_io_service.reset();
 	m_socket.close();
 	m_authorized = false;
 	m_connected = false;
+	
 
+	
 	if (!m_failover.host.empty())
 	{
 		m_retries++;
@@ -134,10 +143,12 @@ void EthStratumClient::resolve_handler(const boost::system::error_code& ec, tcp:
 
 void EthStratumClient::connect_handler(const boost::system::error_code& ec, tcp::resolver::iterator i)
 {
+	dev::setThreadName("stratum");
+	
 	if (!ec)
 	{
 		m_connected = true;
-		cnote << "Connected to stratum server " << p_active->host << ":" << p_active->port;
+		cnote << "Connected to stratum server " << i->host_name() << ":" << p_active->port;
 		if (!p_farm->isMining())
 		{
 			cnote << "Starting farm";
@@ -165,7 +176,7 @@ void EthStratumClient::connect_handler(const boost::system::error_code& ec, tcp:
 }
 
 void EthStratumClient::readline() {
-	m_mtx.lock();
+	x_pending.lock();
 	if (m_pending == 0) {
 		async_read_until(m_socket, m_responseBuffer, "\n",
 			boost::bind(&EthStratumClient::readResponse, this,
@@ -174,7 +185,7 @@ void EthStratumClient::readline() {
 		m_pending++;
 		
 	}
-	m_mtx.unlock();
+	x_pending.unlock();
 }
 
 void EthStratumClient::handleResponse(const boost::system::error_code& ec) {
@@ -184,17 +195,19 @@ void EthStratumClient::handleResponse(const boost::system::error_code& ec) {
 	}
 	else
 	{
+		dev::setThreadName("stratum");
 		cwarn << "Handle response failed: " << ec.message();
 	}
 }
 
 void EthStratumClient::readResponse(const boost::system::error_code& ec, std::size_t bytes_transferred)
 {
-	m_mtx.lock();
+	dev::setThreadName("stratum");
+	x_pending.lock();
 	m_pending = m_pending > 0 ? m_pending - 1 : 0;
-	m_mtx.unlock();
+	x_pending.unlock();
 
-	if (!ec)
+	if (!ec && bytes_transferred)
 	{
 		std::istream is(&m_responseBuffer);
 		std::string response;
@@ -218,12 +231,14 @@ void EthStratumClient::readResponse(const boost::system::error_code& ec, std::si
 		{
 			cwarn << "Discarding incomplete response";
 		}
-		readline();
+		if (m_connected)
+			readline();
 	}
 	else
 	{
 		cwarn << "Read response failed: " << ec.message();
-		reconnect();
+		if (m_connected)
+			reconnect();
 	}
 }
 
@@ -253,16 +268,21 @@ void EthStratumClient::processReponse(Json::Value& responseObject)
 		m_authorized = responseObject.get("result", Json::Value::null).asBool();
 		if (!m_authorized)
 		{
+			cnote << "Worker not authorized:" << p_active->user;
 			disconnect();
 			return;
 		}
 		cnote << "Authorized worker " << p_active->user;
 		break;
 	case 4:
- 		if (responseObject.get("result", false).asBool())
+		if (responseObject.get("result", false).asBool()) {
 			cnote << "B-) Submitted and accepted.";
-		else
+			p_farm->acceptedSolution(m_stale);
+		}
+		else {
 			cwarn << ":-( Not accepted.";
+			p_farm->rejectedSolution(m_stale);
+		}
 		break;
 	default:
 		string method = responseObject.get("method", "").asString();
@@ -271,23 +291,24 @@ void EthStratumClient::processReponse(Json::Value& responseObject)
 			params = responseObject.get("params", Json::Value::null);
 			if (params.isArray())
 			{
-				m_job = params.get((Json::Value::ArrayIndex)0, "").asString();
+				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();
-
-				if (sShareTarget.length() < 66)
-					sShareTarget = "0x" + string(66 - sShareTarget.length(), '0') + sShareTarget.substr(2);
-
+				//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 #" + m_job;
-					cnote << "Header hash: " + sHeaderHash;
-					cnote << "Seed hash: " + sSeedHash;
-					cnote << "Share target: " + 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);
@@ -295,17 +316,38 @@ void EthStratumClient::processReponse(Json::Value& responseObject)
 
 
 					if (seedHash != m_current.seedHash)
+					{
 						cnote << "Grabbing DAG for" << seedHash;
-					if (!(dag = EthashAux::full(seedHash, true, [&](unsigned _pc){ cout << "\rCreating DAG. " << _pc << "% done..." << flush; return 0; })))
+					}
+					if (!(dag = EthashAux::full(seedHash, true, [&](unsigned _pc){ m_waitState = _pc < 100 ? MINER_WAIT_STATE_DAG : MINER_WAIT_STATE_WORK;  cnote << "Creating DAG. " << _pc << "% done..."; return 0; })))
+					{
 						BOOST_THROW_EXCEPTION(DAGCreationFailure());
+					}
 					if (m_precompute)
+					{
 						EthashAux::computeFull(sha3(seedHash), true);
+					}
 					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));
+
 					}
 				}
 			}
@@ -326,29 +368,53 @@ void EthStratumClient::processReponse(Json::Value& responseObject)
 
 }
 
+void EthStratumClient::work_timeout_handler(const boost::system::error_code& ec) {
+	if (!ec) {
+		cnote << "No new work received in" << m_worktimeout << "seconds.";
+		reconnect();
+	}
+}
+
 bool EthStratumClient::submit(EthashProofOfWork::Solution solution) {
+	x_current.lock();
+	EthashProofOfWork::WorkPackage tempWork(m_current);
+	string temp_job = m_job;
+	EthashProofOfWork::WorkPackage tempPreviousWork(m_previous);
+	string temp_previous_job = m_previousJob;
+	x_current.unlock();
 
 	cnote << "Solution found; Submitting to" << p_active->host << "...";
-	cnote << "  Nonce:" << "0x"+solution.nonce.hex();
-	cnote << "  Mixhash:" << "0x" + solution.mixHash.hex();
-	cnote << "  Header-hash:" << "0x" + m_current.headerHash.hex();
-	cnote << "  Seedhash:" << "0x" + m_current.seedHash.hex();
-	cnote << "  Target: " << "0x" + h256(m_current.boundary).hex();
-	cnote << "  Ethash: " << "0x" + h256(EthashAux::eval(m_current.seedHash, m_current.headerHash, solution.nonce).value).hex();
-	if (EthashAux::eval(m_current.seedHash, m_current.headerHash, solution.nonce).value < m_current.boundary)
+	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 + "\",\"" + m_job + "\",\"0x" + solution.nonce.hex() + "\",\"0x" + m_current.headerHash.hex() + "\",\"0x" + solution.mixHash.hex() + "\"]}\n";
+		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";
 		std::ostream os(&m_requestBuffer);
 		os << json;
-
+		m_stale = false;
 		async_write(m_socket, m_requestBuffer,
 			boost::bind(&EthStratumClient::handleResponse, this,
 			boost::asio::placeholders::error));
 		return true;
 	}
-	else
+	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";
+		std::ostream os(&m_requestBuffer);
+		os << json;
+		m_stale = true;
+		cwarn << "Submitting stale solution.";
+		async_write(m_socket, m_requestBuffer,
+			boost::bind(&EthStratumClient::handleResponse, this,
+			boost::asio::placeholders::error));
+		return true;
+	}
+	else {
+		m_stale = false;
 		cwarn << "FAILURE: GPU gave incorrect result!";
-	
+		p_farm->failedSolution();
+	}
+
 	return false;
 }
 

libstratum/EthStratumClient.h

@@ -27,7 +27,7 @@ typedef struct {
 class EthStratumClient
 {
 public:
-	EthStratumClient(GenericFarm<EthashProofOfWork> * f, MinerType m, string const & host, string const & port, string const & user, string const & pass, int const & retries, bool const & precompute);
+	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 const & precompute);
 	~EthStratumClient();
 
 	void setFailover(string const & host, string const & port);
@@ -37,14 +37,17 @@ public:
 	bool isConnected() { return m_connected; }
 	h256 currentHeaderHash() { return m_current.headerHash; }
 	bool current() { return m_current; }
+	unsigned waitState() { return m_waitState; }
 	bool submit(EthashProofOfWork::Solution solution);
+	void reconnect();
 private:
 	void connect();
-	void reconnect();
+	
 	void disconnect();
 	void resolve_handler(const boost::system::error_code& ec, tcp::resolver::iterator i);
 	void connect_handler(const boost::system::error_code& ec, tcp::resolver::iterator i);
-	
+	void work_timeout_handler(const boost::system::error_code& ec);
+
 	void readline();
 	void handleResponse(const boost::system::error_code& ec);
 	void readResponse(const boost::system::error_code& ec, std::size_t bytes_transferred);
@@ -63,14 +66,23 @@ private:
 
 	int	m_retries = 0;
 	int	m_maxRetries;
+	int m_worktimeout = 60;
 
-	boost::mutex m_mtx;
+	int m_waitState = MINER_WAIT_STATE_WORK;
+
+	boost::mutex x_pending;
 	int m_pending;
 	string m_response;
 
 	GenericFarm<EthashProofOfWork> * p_farm;
+	boost::mutex x_current;
 	EthashProofOfWork::WorkPackage m_current;
+	EthashProofOfWork::WorkPackage m_previous;
+
+	bool m_stale = false;
+
 	string m_job;
+	string m_previousJob;
 	EthashAux::FullType m_dag;
 
 	boost::asio::io_service m_io_service;
@@ -78,4 +90,7 @@ private:
 
 	boost::asio::streambuf m_requestBuffer;
 	boost::asio::streambuf m_responseBuffer;
+
+	boost::asio::deadline_timer * p_worktimer;
+
 };