NiceHash modifications

- Added support for EthereumStratum/1.0.0
- Added speed&DAGprogress reporting

Author: nicehashdev

BuildInfo.h.in

@@ -1,6 +1,6 @@
 #pragma once
 
-#define ETH_PROJECT_VERSION "@PROJECT_VERSION@-genoil-@GENOIL_VERSION@"
+#define ETH_PROJECT_VERSION "@PROJECT_VERSION@-genoil-@GENOIL_VERSION@-NiceHash-@NICEHASH_VERSION@"
 #define ETH_COMMIT_HASH @ETH_COMMIT_HASH@
 #define ETH_CLEAN_REPO @ETH_CLEAN_REPO@
 #define ETH_BUILD_TYPE @ETH_BUILD_TYPE@

CMakeLists.txt

@@ -3,6 +3,7 @@ cmake_minimum_required(VERSION 2.8.12)
 
 set(PROJECT_VERSION "0.9.41")
 set(GENOIL_VERSION "1.0.8")
+set(NICEHASH_VERSION "1.0.0")
 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()
@@ -172,7 +173,7 @@ function(createBuildInfo)
 		WORKING_DIRECTORY ${CMAKE_SOURCE_DIR}
 		COMMAND ${CMAKE_COMMAND} -DETH_SOURCE_DIR="${CMAKE_SOURCE_DIR}" -DETH_DST_DIR="${CMAKE_BINARY_DIR}"
 			-DETH_BUILD_TYPE="${_cmake_build_type}" -DETH_BUILD_PLATFORM="${ETH_BUILD_PLATFORM}"
-			-DPROJECT_VERSION="${PROJECT_VERSION}" -DGENOIL_VERSION="${GENOIL_VERSION}" -DETH_FATDB="${FATDB}"
+			-DPROJECT_VERSION="${PROJECT_VERSION}" -DGENOIL_VERSION="${GENOIL_VERSION}" -DNICEHASH_VERSION="${NICEHASH_VERSION}" -DETH_FATDB="${FATDB}"
 			-P "${ETH_SCRIPTS_DIR}/buildinfo.cmake"
 		)
 	include_directories(${CMAKE_CURRENT_BINARY_DIR})

ethminer/MinerAux.h

@@ -96,6 +96,12 @@ struct MiningChannel: public LogChannel
 };
 #define minelog clog(MiningChannel)
 
+struct ReportStruct
+{
+	double speed;
+	unsigned int DAGprogress;
+};
+
 class MinerCLI
 {
 public:
@@ -109,7 +115,12 @@ class MinerCLI
 		Stratum
 	};
 
-	MinerCLI(OperationMode _mode = OperationMode::None): mode(_mode) {}
+	MinerCLI(OperationMode _mode = OperationMode::None): mode(_mode) 
+	{
+		m_speedReportSocket = new boost::asio::ip::udp::socket(m_io_service, boost::asio::ip::udp::v4());
+		boost::asio::socket_base::non_blocking_io command(true);
+		m_speedReportSocket->io_control(command);
+	}
 
 	bool interpretOption(int& i, int argc, char** argv)
 	{
@@ -223,6 +234,10 @@ class MinerCLI
 		{
 			m_worktimeout = atoi(argv[++i]);
 		}
+		else if ((arg == "--report-port") && i + 1 < argc)
+		{
+			m_speedReportPort = atoi(argv[++i]);
+		}
 
 #endif
 #if ETH_ETHASHCL || !ETH_TRUE
@@ -352,8 +367,8 @@ class MinerCLI
 				cerr << "Bad " << arg << " option: " << argv[i] << endl;
 				BOOST_THROW_EXCEPTION(BadArgument());
 			}
-		else if (arg == "-C" || arg == "--cpu")
-			m_minerType = MinerType::CPU;
+		//else if (arg == "-C" || arg == "--cpu")
+		//	m_minerType = MinerType::CPU;
 		else if (arg == "-G" || arg == "--opencl")
 			m_minerType = MinerType::CL;
 		else if (arg == "-U" || arg == "--cuda")
@@ -624,7 +639,7 @@ class MinerCLI
 			<< "		bench - like old, but keep epoch 0 for benchmarking" << endl
 			<< "        all   - erase all DAG files. After deleting all files, setting changes to none." << endl
 			<< "Mining configuration:" << endl
-			<< "    -C,--cpu  When mining, use the CPU." << endl
+			//<< "    -C,--cpu  When mining, use the CPU." << endl
 			<< "    -G,--opencl  When mining use the GPU via OpenCL." << endl
 			<< "    -U,--cuda  When mining use the GPU via CUDA." << endl
 			<< "    --opencl-platform <n>  When mining using -G/--opencl use OpenCL platform n (default: 0)." << endl
@@ -651,6 +666,7 @@ class MinerCLI
 			<< "        sync  - Instruct CUDA to block the CPU thread on a synchronization primitive when waiting for the results from the device." << endl
 			<< "    --cuda-devices <0 1 ..n> Select which CUDA GPUs to mine on. Default is to use all" << endl
 #endif
+			<< "    --report-port Speed reporting port (used by NiceHash Miner)" << endl
 			;
 	}
 
@@ -662,11 +678,22 @@ class MinerCLI
 	{
 		h256 seedHash = EthashAux::seedHash(_n);
 		cout << "Initializing DAG for epoch beginning #" << (_n / 30000 * 30000) << " (seedhash " << seedHash.abridged() << "). This will take a while." << endl;
-		EthashAux::full(seedHash, true);
+		EthashAux::full(seedHash, true, [&](unsigned _pc) {
+			cout << "\rCreating DAG. " << _pc << "% done..." << flush;
+			reportDAGprogress(_pc);
+			return 0;
+		});
 		exit(0);
 	}
 
-	
+	void reportDAGprogress(unsigned _pc)
+	{
+		ReportStruct rs;
+		rs.DAGprogress = _pc;
+		rs.speed = 0;
+		boost::asio::ip::udp::endpoint endpoint(boost::asio::ip::address::from_string("127.0.0.1"), m_speedReportPort);
+		m_speedReportSocket->send_to(boost::asio::buffer((void*)&rs, (size_t)sizeof(ReportStruct)), endpoint);
+	}
 
 	void doBenchmark(MinerType _m, bool _phoneHome, unsigned _warmupDuration = 15, unsigned _trialDuration = 3, unsigned _trials = 5)
 	{
@@ -1049,23 +1076,64 @@ class MinerCLI
 			client.submit(sol);
 			return false;
 		});
+
+		boost::asio::ip::udp::endpoint endpoint(boost::asio::ip::address::from_string("127.0.0.1"), m_speedReportPort);
 
+#define WORKINGPROGRESS_BACKLOG_SIZE 32
+#define REPORT_DELAY 4
+		WorkingProgress wplist[WORKINGPROGRESS_BACKLOG_SIZE];
+		int wplist_index = 0;
+		bool first = false;
+		//int wplist_filled = 0;
+
 		while (client.isRunning())
 		{
 			auto mp = f.miningProgress();
+			if (!first) first = true; // for some reason, first ms is huge...
+			else
+			{
+				wplist[wplist_index].hashes = mp.hashes;
+				wplist[wplist_index].ms = mp.ms;
+				wplist_index = (wplist_index + 1) % WORKINGPROGRESS_BACKLOG_SIZE;
+				//wplist_filled = (wplist_filled == WORKINGPROGRESS_BACKLOG_SIZE) ? wplist_filled : (wplist_filled + 1);
+				mp.hashes = 0;
+				mp.ms = 0;
+				for (int i = 0; i != WORKINGPROGRESS_BACKLOG_SIZE; ++i)
+				{
+					mp.hashes += wplist[i].hashes;
+					mp.ms += wplist[i].ms;
+				}
+
+				if (wplist_index % REPORT_DELAY == 0)
+				{
+					ReportStruct rs;
+					rs.speed = 0;
+					if (mp.ms > 0)
+						rs.speed = (double)mp.hashes / (mp.ms * 1000);
+					rs.DAGprogress = 100;
+					m_speedReportSocket->send_to(boost::asio::buffer((void*)&rs, sizeof(ReportStruct)), endpoint);
+				}
+			}
+
 			f.resetMiningProgress();
 			if (client.isConnected())
 			{
 				if (client.current())
-					minelog << "Mining on PoWhash" << "#"+(client.currentHeaderHash().hex().substr(0,8)) << ": " << mp << f.getSolutionStats();
+					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));
 		}
+
+		delete m_speedReportSocket;
 	}
 #endif
 
+	uint16_t m_speedReportPort = 37450;
+	boost::asio::ip::udp::socket* m_speedReportSocket;
+	boost::asio::io_service m_io_service;
+
 	/// Operating mode.
 	OperationMode mode;
 	DAGEraseMode m_eraseMode = DAGEraseMode::None;

ethminer/main.cpp

@@ -70,6 +70,7 @@ int main(int argc, char** argv)
 {
 	cout << "Genoil's ethminer " << ETH_PROJECT_VERSION << endl;
 	cout << "=====================================================================" << endl;
+	cout << "Modified by NiceHash with better&correct stratum" << endl;
 	cout << "Forked from github.com/ethereum/cpp-ethereum" << endl;
 	cout << "CUDA kernel ported from Tim Hughes' OpenCL kernel" << endl;
 	cout << "With contributions from nerdralph, RoBiK, tpruvot and sp_ " << endl << endl;

libethash-cl/ethash_cl_miner.cpp

@@ -471,7 +471,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, uint64_t startn)
 {
 	try
 	{
@@ -497,8 +497,9 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook
 		m_searchKernel.setArg(4, target);
 
 		unsigned buf = 0;
-		random_device engine;
-		uint64_t start_nonce = uniform_int_distribution<uint64_t>()(engine);
+		//random_device engine;
+		//uint64_t start_nonce = uniform_int_distribution<uint64_t>()(engine);
+		uint64_t start_nonce = startn;
 		for (;; start_nonce += m_globalWorkSize)
 		{
 			// supply output buffer to kernel

libethash-cl/ethash_cl_miner.h

@@ -58,7 +58,7 @@ class ethash_cl_miner
 		unsigned _deviceId = 0
 	);
 	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, uint64_t startn);
 
 	/* -- default values -- */
 	/// Default value of the local work size. Also known as workgroup size.

libethash-cuda/ethash_cuda_miner.cpp

@@ -251,7 +251,7 @@ bool ethash_cuda_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned _d
 	}
 }
 
-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, uint64_t startn)
 {
 	bool initialize = false;
 	bool exit = false;
@@ -269,13 +269,20 @@ void ethash_cuda_miner::search(uint8_t const* header, uint64_t target, search_ho
 	}
 	if (initialize)
 	{
-		random_device engine;
-		m_current_nonce = uniform_int_distribution<uint64_t>()(engine);
+		//random_device engine;
+		//m_current_nonce = uniform_int_distribution<uint64_t>()(engine);
+		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;
+	}
 	uint64_t batch_size = s_gridSize * s_blockSize;
 	for (; !exit; m_current_index++, m_current_nonce += batch_size)
 	{

libethash-cuda/ethash_cuda_miner.h

@@ -40,7 +40,7 @@ class ethash_cuda_miner
 		unsigned _deviceId = 0
 		);
 	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, uint64_t startn);
 
 	/* -- default values -- */
 	/// Default value of the block size. Also known as workgroup size.
@@ -54,6 +54,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;
 
 	volatile uint32_t ** m_search_buf;

libethcore/EthashAux.h

@@ -63,6 +63,8 @@ struct EthashProofOfWork
 		h256 boundary;
 		h256 headerHash;	///< When h256() means "pause until notified a new work package is available".
 		h256 seedHash;
+		uint64_t startNonce;
+		int exSizeBits;
 	};
 
 	static const WorkPackage NullWorkPackage;

libethcore/EthashCUDAMiner.cpp

@@ -31,6 +31,7 @@ along with cpp-ethereum.  If not, see <http://www.gnu.org/licenses/>.
 #include <thread>
 #include <chrono>
 #include <libethash-cuda/ethash_cuda_miner.h>
+#include <libethash/endian.h>
 using namespace std;
 using namespace dev;
 using namespace eth;
@@ -180,7 +181,11 @@ void EthashCUDAMiner::workLoop()
 		}
 
 		uint64_t upper64OfBoundary = (uint64_t)(u64)((u256)w.boundary >> 192);
-		m_miner->search(w.headerHash.data(), upper64OfBoundary, *m_hook);
+		uint64_t startn = w.startNonce | ((uint64_t)index() << (64 - 4 - w.exSizeBits)); // this can support up to 16 devices
+		uint64_t swapped = ethash_swap_u64(startn);
+		Nonce startN((byte const*)&swapped, h64::ConstructFromPointerType::ConstructFromPointer);
+		cnote << "starting nonce is " << startN.hex();
+		m_miner->search(w.headerHash.data(), upper64OfBoundary, *m_hook, startn);
 	}
 	catch (std::runtime_error const& _e)
 	{