diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp
index 667144cdd3a2ea060c1d1ca395e88d07b6ada42d..0a181154c381ccb1fbe1dbc3d91f66888c9f4a94 100644
--- a/xmrstak/backend/amd/minethd.cpp
+++ b/xmrstak/backend/amd/minethd.cpp
@@ -172,7 +172,6 @@ void minethd::work_main()
 	lck.release();
 	std::this_thread::yield();
 
-	uint64_t iCount = 0;
 	cryptonight_ctx* cpu_ctx;
 	cpu_ctx = cpu::minethd::minethd_alloc_ctx();
 
@@ -288,10 +287,7 @@ void minethd::work_main()
 					executor::inst()->push_event(ex_event("AMD Invalid Result", pGpuCtx->deviceIdx, oWork.iPoolId));
 			}
 
-			iCount += pGpuCtx->rawIntensity;
-			uint64_t iStamp = get_timestamp_ms();
-			iHashCount.store(iCount, std::memory_order_relaxed);
-			iTimestamp.store(iStamp, std::memory_order_relaxed);
+			updateStats(pGpuCtx->rawIntensity, oWork.iPoolId);
 
 			accRuntime += updateTimings(pGpuCtx, t0);
 
diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp
index 0229af0a75c702b6ae823ad91ccedfd68fef1441..463be1aabea509de74784912ae05e1e452456035 100644
--- a/xmrstak/backend/cpu/minethd.cpp
+++ b/xmrstak/backend/cpu/minethd.cpp
@@ -833,6 +833,7 @@ void minethd::multiway_work_main()
 
 	cryptonight_ctx* ctx[MAX_N];
 	uint64_t iCount = 0;
+	uint64_t iLastCount = 0;
 	uint64_t* piHashVal[MAX_N];
 	uint32_t* piNonce[MAX_N];
 	uint8_t bHashOut[MAX_N * 32];
@@ -915,9 +916,8 @@ void minethd::multiway_work_main()
 		{
 			if((iCount++ & 0x7) == 0) //Store stats every 8*N hashes
 			{
-				uint64_t iStamp = get_timestamp_ms();
-				iHashCount.store(iCount * N, std::memory_order_relaxed);
-				iTimestamp.store(iStamp, std::memory_order_relaxed);
+				updateStats((iCount - iLastCount) * N, oWork.iPoolId);
+				iLastCount = iCount;
 			}
 
 			nonce_ctr -= N;
diff --git a/xmrstak/backend/iBackend.hpp b/xmrstak/backend/iBackend.hpp
index 1af42c248964accea4f9ead8eb8e703341b378fa..dd59b6c52e906a7da61402dbc86c842b91f3ca43 100644
--- a/xmrstak/backend/iBackend.hpp
+++ b/xmrstak/backend/iBackend.hpp
@@ -1,6 +1,7 @@
 #pragma once
 
 #include "xmrstak/backend/globalStates.hpp"
+#include "xmrstak/net/msgstruct.hpp"
 
 #include <atomic>
 #include <climits>
@@ -46,6 +47,29 @@ struct iBackend
 	std::atomic<uint64_t> iTimestamp;
 	uint32_t iThreadNo;
 	BackendType backendType = UNKNOWN;
+	uint64_t iLastStamp = get_timestamp_ms();
+	double avgHashPerMsec = 0.0;
+
+	void updateStats(uint64_t numNewHashes, size_t poolId)
+	{
+		uint64_t iStamp = get_timestamp_ms();
+		double timeDiff = static_cast<double>(iStamp - iLastStamp);
+		iLastStamp = iStamp;
+
+		if(poolId == 0)
+		{
+			// if dev pool is active interpolate the number of shares (avoid hash rate drops)
+			numNewHashes = static_cast<uint64_t>(avgHashPerMsec * timeDiff);
+		}
+		else
+		{
+			const double hashRatePerMs = static_cast<double>(numNewHashes) / timeDiff;
+			constexpr double averagingBias = 0.1;
+			avgHashPerMsec = avgHashPerMsec * (1.0 - averagingBias) + hashRatePerMs * averagingBias;
+		}
+		iHashCount.fetch_add(numNewHashes, std::memory_order_relaxed);
+		iTimestamp.store(iStamp, std::memory_order_relaxed);
+	}
 
 	iBackend() :
 		iHashCount(0),
diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp
index 4506faed6d39fd3554607c5eca9a2ff6bf4f420b..32b21dc7105079ccad14e67baedea9772f5d062d 100644
--- a/xmrstak/backend/nvidia/minethd.cpp
+++ b/xmrstak/backend/nvidia/minethd.cpp
@@ -198,7 +198,6 @@ void minethd::work_main()
 	// wait until all NVIDIA devices are initialized
 	thread_work_guard.wait();
 
-	uint64_t iCount = 0;
 	cryptonight_ctx* cpu_ctx;
 	cpu_ctx = cpu::minethd::minethd_alloc_ctx();
 
@@ -297,13 +296,8 @@ void minethd::work_main()
 					executor::inst()->push_event(ex_event("NVIDIA Invalid Result", ctx.device_id, oWork.iPoolId));
 			}
 
-			iCount += h_per_round;
 			iNonce += h_per_round;
-
-			using namespace std::chrono;
-			uint64_t iStamp = get_timestamp_ms();
-			iHashCount.store(iCount, std::memory_order_relaxed);
-			iTimestamp.store(iStamp, std::memory_order_relaxed);
+			updateStats(h_per_round, oWork.iPoolId);
 			std::this_thread::yield();
 		}
 
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_cryptonight_r.curt b/xmrstak/backend/nvidia/nvcc_code/cuda_cryptonight_r.curt
index bcf49508009f691fdafbdf48f14690c91ab501e2..214114c7ec0a83dd166800d6024dc9006c185b16 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_cryptonight_r.curt
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_cryptonight_r.curt
@@ -462,10 +462,10 @@ __global__ void CryptonightR_phase2(
     uint64_t bx0             = ((uint64_t*)(d_ctx_b + thread * 16))[sub];
     uint64_t bx1             = ((uint64_t*)(d_ctx_b + thread * 16 + 4))[sub];
 
-    uint32_t r0 = d_ctx_b[thread * 16 + 4 * 2];
-    uint32_t r1 = d_ctx_b[thread * 16 + 4 * 2 + 1];
-    uint32_t r2 = d_ctx_b[thread * 16 + 4 * 2 + 2];
-    uint32_t r3 = d_ctx_b[thread * 16 + 4 * 2 + 3];
+    volatile uint32_t r0 = d_ctx_b[thread * 16 + 4 * 2];
+    volatile uint32_t r1 = d_ctx_b[thread * 16 + 4 * 2 + 1];
+    volatile uint32_t r2 = d_ctx_b[thread * 16 + 4 * 2 + 2];
+    volatile uint32_t r3 = d_ctx_b[thread * 16 + 4 * 2 + 3];
 
     const int batchsize      = (ITERATIONS * 2) >> ( 1 + bfactor );
     const int start          = partidx * batchsize;
diff --git a/xmrstak/cli/cli-miner.cpp b/xmrstak/cli/cli-miner.cpp
index a53a4ba9b6f51e94bca223a25d5726851f903718..cc4b1b3d96db5b10ca2330c56abfc4101b1d44ed 100644
--- a/xmrstak/cli/cli-miner.cpp
+++ b/xmrstak/cli/cli-miner.cpp
@@ -195,7 +195,14 @@ inline void prompt_once(bool& prompted)
 	}
 }
 
-void do_guided_pool_config(const bool use_simple_start)
+inline bool use_simple_start()
+{
+	// ask this question only once
+	static bool simple_start = read_yes_no("\nUse simple setup method? (Y/n)", "Y");
+	return simple_start;
+}
+
+void do_guided_pool_config()
 {
 	using namespace xmrstak;
 
@@ -261,19 +268,22 @@ void do_guided_pool_config(const bool use_simple_start)
 	}
 
 	auto& rigid = params::inst().poolRigid;
-	if(!use_simple_start && rigid.empty() && !params::inst().userSetRigid)
+	if(rigid.empty() && !params::inst().userSetRigid)
 	{
-		prompt_once(prompted);
-
-		if(!stdin_flushed)
+		if(!use_simple_start())
 		{
-			// clear everything from stdin to allow an empty rigid
-			std::cin.clear();
-			std::cin.ignore(INT_MAX, '\n');
-		}
+			prompt_once(prompted);
 
-		std::cout << "- Rig identifier for pool-side statistics (needs pool support). Can be empty:" << std::endl;
-		getline(std::cin, rigid);
+			if(!stdin_flushed)
+			{
+				// clear everything from stdin to allow an empty rigid
+				std::cin.clear();
+				std::cin.ignore(INT_MAX, '\n');
+			}
+
+			std::cout << "- Rig identifier for pool-side statistics (needs pool support). Can be empty:" << std::endl;
+			getline(std::cin, rigid);
+		}
 	}
 
 	bool tls = params::inst().poolUseTls;
@@ -289,15 +299,19 @@ void do_guided_pool_config(const bool use_simple_start)
 #endif
 
 	bool nicehash = params::inst().nicehashMode;
-	if(!use_simple_start && !userSetPool)
+	if(!userSetPool)
 	{
-		prompt_once(prompted);
-		nicehash = read_yes_no("- Do you want to use nicehash on this pool? (y/N)", "N");
+		if(!use_simple_start())
+		{
+			prompt_once(prompted);
+			nicehash = read_yes_no("- Do you want to use nicehash on this pool? (y/N)", "N");
+		}
 	}
 
 	bool multipool = false;
-	if(!use_simple_start && !userSetPool)
-		multipool = read_yes_no("- Do you want to use multiple pools? (y/N)", "N");
+	if(!userSetPool)
+		if(!use_simple_start())
+			multipool = read_yes_no("- Do you want to use multiple pools? (y/N)", "N");
 
 	int64_t pool_weight = 1;
 	if(multipool)
@@ -335,7 +349,7 @@ void do_guided_pool_config(const bool use_simple_start)
 	std::cout << "Pool configuration stored in file '" << params::inst().configFilePools << "'" << std::endl;
 }
 
-void do_guided_config(const bool use_simple_start)
+void do_guided_config()
 {
 	using namespace xmrstak;
 
@@ -353,7 +367,7 @@ void do_guided_config(const bool use_simple_start)
 	{
 		http_port = params::httpd_port_disabled;
 #ifndef CONF_NO_HTTPD
-		if(!use_simple_start)
+		if(!use_simple_start())
 		{
 			prompt_once(prompted);
 
@@ -737,17 +751,13 @@ int main(int argc, char* argv[])
 	bool hasConfigFile = configEditor::file_exist(params::inst().configFile);
 	bool hasPoolConfig = configEditor::file_exist(params::inst().configFilePools);
 
-	if(!hasConfigFile || !hasPoolConfig)
-	{
-		bool use_simple_start = read_yes_no("\nUse simple setup method? (Y/n)", "Y");
+	// check if we need a guided start
+	if(!hasConfigFile)
+		do_guided_config();
 
-		// check if we need a guided start
-		if(!hasConfigFile)
-			do_guided_config(use_simple_start);
+	if(!hasPoolConfig)
+		do_guided_pool_config();
 
-		if(!hasPoolConfig)
-			do_guided_pool_config(use_simple_start);
-	}
 	if(!jconf::inst()->parse_config(params::inst().configFile.c_str(), params::inst().configFilePools.c_str()))
 	{
 		win_exit();
diff --git a/xmrstak/version.cpp b/xmrstak/version.cpp
index d4afd7877a0ac52f88b097b5d6f78ee0d5766530..f42e2a0bf0bfe68dc3c3d90f98ef635af6fef03a 100644
--- a/xmrstak/version.cpp
+++ b/xmrstak/version.cpp
@@ -20,7 +20,7 @@
 #endif
 
 #define XMR_STAK_NAME "xmr-stak"
-#define XMR_STAK_VERSION "2.10.3"
+#define XMR_STAK_VERSION "2.10.4"
 
 #if defined(_WIN32)
 #define OS_TYPE "win"