shorter work mutex lock

longer work timeout
binary

ethminer/MinerAux.h

@@ -1115,7 +1115,7 @@ private:
 	unsigned m_farmRecheckPeriod = 500;
 	unsigned m_defaultStratumFarmRecheckPeriod = 2000;
 	bool m_farmRecheckSet = false;
-	int m_worktimeout = 60;
+	int m_worktimeout = 90;
 	bool m_precompute = true;
 
 #if ETH_STRATUM || !ETH_TRUE

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)

libstratum/EthStratumClient.cpp

@@ -329,7 +329,7 @@ void EthStratumClient::processReponse(Json::Value& responseObject)
 					}
 					if (headerHash != m_current.headerHash)
 					{
-						x_current.lock();
+						//x_current.lock();
 						if (p_worktimer)
 							p_worktimer->cancel();
 
@@ -344,7 +344,7 @@ void EthStratumClient::processReponse(Json::Value& responseObject)
 						m_job = job;
 
 						p_farm->setWork(m_current);
-						x_current.unlock();
+						//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));
 
@@ -377,36 +377,32 @@ void EthStratumClient::work_timeout_handler(const boost::system::error_code& ec)
 
 bool EthStratumClient::submit(EthashProofOfWork::Solution solution) {
 	x_current.lock();
-	x_stale.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;
-		x_stale.unlock();
-		x_current.unlock();
 		async_write(m_socket, m_requestBuffer,
 			boost::bind(&EthStratumClient::handleResponse, this,
 			boost::asio::placeholders::error));
 		return true;
 	}
-	else if (EthashAux::eval(m_previous.seedHash, m_previous.headerHash, solution.nonce).value < m_previous.boundary)
+	else if (EthashAux::eval(tempPreviousWork.seedHash, tempPreviousWork.headerHash, solution.nonce).value < tempPreviousWork.boundary)
 	{
-		string json = "{\"id\": 4, \"method\": \"mining.submit\", \"params\": [\"" + p_active->user + "\",\"" + m_previousJob + "\",\"0x" + solution.nonce.hex() + "\",\"0x" + m_previous.headerHash.hex() + "\",\"0x" + solution.mixHash.hex() + "\"]}\n";
+		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;
-		x_stale.unlock();
-		x_current.unlock();
 		cwarn << "Submitting stale solution.";
 		async_write(m_socket, m_requestBuffer,
 			boost::bind(&EthStratumClient::handleResponse, this,
@@ -415,12 +411,10 @@ bool EthStratumClient::submit(EthashProofOfWork::Solution solution) {
 	}
 	else {
 		m_stale = false;
-		x_stale.unlock();
-		x_current.unlock();
 		cwarn << "FAILURE: GPU gave incorrect result!";
 		p_farm->failedSolution();
 	}
-	
+
 	return false;
 }
 

libstratum/EthStratumClient.h

@@ -79,7 +79,6 @@ private:
 	EthashProofOfWork::WorkPackage m_current;
 	EthashProofOfWork::WorkPackage m_previous;
 
-	boost::mutex x_stale;
 	bool m_stale = false;
 
 	string m_job;