From 68ac9673ec09cbb4f32d614348270f245abea30e Mon Sep 17 00:00:00 2001
From: psychocrypt <psychocryptHPC@gmail.com>
Date: Sun, 1 Apr 2018 21:37:44 +0200
Subject: [PATCH] refactor mining algo selection

- add `fork_height` to currency
- refactor algorithm selection
---
 xmrstak/backend/amd/amd_gpu/gpu.cpp           |  6 +-
 xmrstak/backend/amd/minethd.cpp               | 29 ++++-----
 xmrstak/backend/cpu/minethd.cpp               | 60 +++++++++++--------
 xmrstak/backend/cpu/minethd.hpp               |  2 +-
 xmrstak/backend/nvidia/minethd.cpp            | 33 ++++------
 .../backend/nvidia/nvcc_code/cryptonight.hpp  |  6 +-
 xmrstak/backend/nvidia/nvcc_code/cuda_core.cu | 19 ++----
 .../backend/nvidia/nvcc_code/cuda_extra.cu    |  8 +--
 xmrstak/jconf.cpp                             | 22 +++----
 xmrstak/jconf.hpp                             |  7 ++-
 10 files changed, 90 insertions(+), 102 deletions(-)

diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp
index 4a82fab..2973db4 100644
--- a/xmrstak/backend/amd/amd_gpu/gpu.cpp
+++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp
@@ -940,7 +940,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
 
 	/// @todo only activate if currency is monero
 	int cn_kernel_offset = 0;
-	if((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon ) && version >= 7)
+	if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon)
 	{
 		cn_kernel_offset = 6;
 	}
@@ -966,7 +966,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
 		return(ERR_OCL_API);
 	}
 
-	if((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon ) && version >= 7)
+	if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon )
 	{
 		// Input
 		if ((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS)
@@ -1134,7 +1134,7 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo,
 	size_t tmpNonce = ctx->Nonce;
 	/// @todo only activate if currency is monero
 	int cn_kernel_offset = 0;
-	if((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version >= 7)
+	if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon)
 	{
 		cn_kernel_offset = 6;
 	}
diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp
index 78774c0..bc14e44 100644
--- a/xmrstak/backend/amd/minethd.cpp
+++ b/xmrstak/backend/amd/minethd.cpp
@@ -191,10 +191,10 @@ void minethd::work_main()
 	uint64_t iCount = 0;
 	cryptonight_ctx* cpu_ctx;
 	cpu_ctx = cpu::minethd::minethd_alloc_ctx();
-	auto miner_algo = ::jconf::inst()->GetMiningAlgo();
 	
 	// start with root algorithm and switch later if fork version is reached
-	cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, ::jconf::inst()->GetMiningAlgoRoot());
+	auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+	cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
 
 	globalStates::inst().iConsumeCnt++;
 
@@ -213,16 +213,18 @@ void minethd::work_main()
 				std::this_thread::sleep_for(std::chrono::milliseconds(100));
 
 			consume_work();
-			uint8_t new_version = oWork.getVersion();
-			const bool time_to_fork =
-				((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version < 7 && new_version >= 7) ||
-				(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3);
-			if(time_to_fork)
+			continue;
+		}
+
+		uint8_t new_version = oWork.getVersion();
+		if(new_version != version)
+		{
+			if(new_version >= ::jconf::inst()->GetMiningForkHeight())
 			{
+				miner_algo = ::jconf::inst()->GetMiningAlgo();
 				hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
 			}
 			version = new_version;
-			continue;
 		}
 
 		uint32_t h_per_round = pGpuCtx->rawIntensity;
@@ -230,7 +232,7 @@ void minethd::work_main()
 
 		assert(sizeof(job_result::sJobID) == sizeof(pool_job::sJobID));
 		uint64_t target = oWork.iTarget;
-		/// \todo add monero hard for version
+		
 		XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target, miner_algo, version);
 
 		if(oWork.bNiceHash)
@@ -274,15 +276,6 @@ void minethd::work_main()
 		}
 
 		consume_work();
-		uint8_t new_version = oWork.getVersion();
-		const bool time_to_fork =
-			((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version < 7 && new_version >= 7) ||
-			(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3);
-		if(time_to_fork)
-		{
-			hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
-		}
-		version = new_version;
 	}
 }
 
diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp
index d7e242b..4c55515 100644
--- a/xmrstak/backend/cpu/minethd.cpp
+++ b/xmrstak/backend/cpu/minethd.cpp
@@ -417,14 +417,15 @@ void minethd::work_main()
 	lck.release();
 	std::this_thread::yield();
 
-	cn_hash_fun hash_fun;
 	cryptonight_ctx* ctx;
 	uint64_t iCount = 0;
 	uint64_t* piHashVal;
 	uint32_t* piNonce;
 	job_result result;
 
-	hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo());
+	// start with root algorithm and switch later if fork version is reached
+	auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+	cn_hash_fun hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
 	ctx = minethd_alloc_ctx();
 
 	piHashVal = (uint64_t*)(result.bResult + 24);
@@ -432,6 +433,8 @@ void minethd::work_main()
 	globalStates::inst().inst().iConsumeCnt++;
 	result.iThreadId = iThreadNo;
 
+	uint8_t version = 0;
+
 	while (bQuit == 0)
 	{
 		if (oWork.bStall)
@@ -457,15 +460,16 @@ void minethd::work_main()
 		if(oWork.bNiceHash)
 			result.iNonce = *piNonce;
 
-		auto miner_algo = ::jconf::inst()->GetMiningAlgo();
-		const bool time_to_fork =
-			((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && oWork.bWorkBlob[0] >= 7) ||
-			(miner_algo == cryptonight_heavy && oWork.bWorkBlob[0] >= 3);
-
-		if(time_to_fork)
-			hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
-		else
-			hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgoRoot());
+		uint8_t new_version = oWork.getVersion();
+		if(new_version != version)
+		{
+			if(new_version >= ::jconf::inst()->GetMiningForkHeight())
+			{
+				miner_algo = ::jconf::inst()->GetMiningAlgo();
+				hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
+			}
+			version = new_version;
+		}
 
 		while(globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo)
 		{
@@ -623,22 +627,22 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
 
 void minethd::double_work_main()
 {
-	multiway_work_main<2>(func_multi_selector(2, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo()));
+	multiway_work_main<2>();
 }
 
 void minethd::triple_work_main()
 {
-	multiway_work_main<3>(func_multi_selector(3, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo()));
+	multiway_work_main<3>();
 }
 
 void minethd::quad_work_main()
 {
-	multiway_work_main<4>(func_multi_selector(4, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo()));
+	multiway_work_main<4>();
 }
 
 void minethd::penta_work_main()
 {
-	multiway_work_main<5>(func_multi_selector(5, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo()));
+	multiway_work_main<5>();
 }
 
 template<size_t N>
@@ -653,7 +657,7 @@ void minethd::prep_multiway_work(uint8_t *bWorkBlob, uint32_t **piNonce)
 }
 
 template<size_t N>
-void minethd::multiway_work_main(cn_hash_fun_multi hash_fun_multi)
+void minethd::multiway_work_main()
 {
 	if(affinity >= 0) //-1 means no affinity
 		bindMemoryToNUMANode(affinity);
@@ -684,6 +688,11 @@ void minethd::multiway_work_main(cn_hash_fun_multi hash_fun_multi)
 
 	globalStates::inst().iConsumeCnt++;
 
+	// start with root algorithm and switch later if fork version is reached
+	auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+	cn_hash_fun_multi hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
+	uint8_t version = 0;
+
 	while (bQuit == 0)
 	{
 		if (oWork.bStall)
@@ -708,15 +717,16 @@ void minethd::multiway_work_main(cn_hash_fun_multi hash_fun_multi)
 		if(oWork.bNiceHash)
 			iNonce = *piNonce[0];
 
-		auto miner_algo = ::jconf::inst()->GetMiningAlgo();
-		const bool time_to_fork =
-			((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && oWork.bWorkBlob[0] >= 7) ||
-			(miner_algo == cryptonight_heavy && oWork.bWorkBlob[0] >= 3);
-
-		if(time_to_fork)
-			hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
-		else
-			hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgoRoot());
+		uint8_t new_version = oWork.getVersion();
+		if(new_version != version)
+		{
+			if(new_version >= ::jconf::inst()->GetMiningForkHeight())
+			{
+				miner_algo = ::jconf::inst()->GetMiningAlgo();
+				hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
+			}
+			version = new_version;
+		}
 
 		while (globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo)
 		{
diff --git a/xmrstak/backend/cpu/minethd.hpp b/xmrstak/backend/cpu/minethd.hpp
index ef1bbd2..cd78343 100644
--- a/xmrstak/backend/cpu/minethd.hpp
+++ b/xmrstak/backend/cpu/minethd.hpp
@@ -36,7 +36,7 @@ private:
 	minethd(miner_work& pWork, size_t iNo, int iMultiway, bool no_prefetch, int64_t affinity);
 
 	template<size_t N>
-	void multiway_work_main(cn_hash_fun_multi hash_fun_multi);
+	void multiway_work_main();
 
 	template<size_t N>
 	void prep_multiway_work(uint8_t *bWorkBlob, uint32_t **piNonce);
diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp
index f8e8409..917101f 100644
--- a/xmrstak/backend/nvidia/minethd.cpp
+++ b/xmrstak/backend/nvidia/minethd.cpp
@@ -237,10 +237,10 @@ void minethd::work_main()
 	uint64_t iCount = 0;
 	cryptonight_ctx* cpu_ctx;
 	cpu_ctx = cpu::minethd::minethd_alloc_ctx();
-	auto miner_algo = ::jconf::inst()->GetMiningAlgo();
 	
 	// start with root algorithm and switch later if fork version is reached
-	cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, ::jconf::inst()->GetMiningAlgoRoot());
+	auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+	cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
 	
 	uint32_t iNonce;
 
@@ -261,16 +261,17 @@ void minethd::work_main()
 				std::this_thread::sleep_for(std::chrono::milliseconds(100));
 
 			consume_work();
-			uint8_t new_version = oWork.getVersion();
-			const bool time_to_fork =
-				((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version < 7 && new_version >= 7) ||
-				(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3);
-			if(time_to_fork)
+			continue;
+		}
+		uint8_t new_version = oWork.getVersion();
+		if(new_version != version)
+		{
+			if(new_version >= ::jconf::inst()->GetMiningForkHeight())
 			{
+				miner_algo = ::jconf::inst()->GetMiningAlgo();
 				hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
 			}
 			version = new_version;
-			continue;
 		}
 
 		cryptonight_extra_cpu_set_data(&ctx, oWork.bWorkBlob, oWork.iWorkSize);
@@ -294,11 +295,11 @@ void minethd::work_main()
 			uint32_t foundNonce[10];
 			uint32_t foundCount;
 
-			cryptonight_extra_cpu_prepare(&ctx, iNonce, miner_algo, version);
+			cryptonight_extra_cpu_prepare(&ctx, iNonce, miner_algo);
 
-			cryptonight_core_cpu_hash(&ctx, miner_algo, iNonce, version);
+			cryptonight_core_cpu_hash(&ctx, miner_algo, iNonce);
 
-			cryptonight_extra_cpu_final(&ctx, iNonce, oWork.iTarget, &foundCount, foundNonce, miner_algo, version);
+			cryptonight_extra_cpu_final(&ctx, iNonce, oWork.iTarget, &foundCount, foundNonce, miner_algo);
 
 			for(size_t i = 0; i < foundCount; i++)
 			{
@@ -329,18 +330,8 @@ void minethd::work_main()
 		}
 
 		consume_work();
-		uint8_t new_version = oWork.getVersion();
-		const bool time_to_fork =
-			((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version < 7 && new_version >= 7) ||
-			(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3);
-		if(time_to_fork)
-		{
-			hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
-		}
-		version = new_version;
 	}
 }
 
 } // namespace xmrstak
-
 } //namespace nvidia
diff --git a/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp b/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp
index 29a3523..c1e31b9 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp
+++ b/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp
@@ -45,8 +45,8 @@ int cuda_get_devicecount( int* deviceCount);
 int cuda_get_deviceinfo(nvid_ctx *ctx);
 int cryptonight_extra_cpu_init(nvid_ctx *ctx);
 void cryptonight_extra_cpu_set_data( nvid_ctx* ctx, const void *data, uint32_t len);
-void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce, xmrstak_algo miner_algo, uint8_t version);
-void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce,xmrstak_algo miner_algo, uint8_t version);
+void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce, xmrstak_algo miner_algo);
+void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce,xmrstak_algo miner_algo);
 }
 
-void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t startNonce, uint8_t version);
+void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t startNonce);
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
index a1a9e15..7aa44e8 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
@@ -475,24 +475,16 @@ void cryptonight_core_gpu_hash(nvid_ctx* ctx, uint32_t nonce)
 	}
 }
 
-void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t startNonce, uint8_t version)
+void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t startNonce)
 {
 
 	if(miner_algo == cryptonight_monero)
 	{
-		if(version >= 7)
-			cryptonight_core_gpu_hash<CRYPTONIGHT_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight_monero>(ctx, startNonce);
-		else
-			cryptonight_core_gpu_hash<CRYPTONIGHT_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight>(ctx, startNonce);
+		cryptonight_core_gpu_hash<CRYPTONIGHT_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight_monero>(ctx, startNonce);
 	}
 	else if(miner_algo == cryptonight_heavy)
 	{
-		if(version >= 3)
-			cryptonight_core_gpu_hash<CRYPTONIGHT_HEAVY_ITER, CRYPTONIGHT_HEAVY_MASK, CRYPTONIGHT_HEAVY_MEMORY/4, cryptonight_heavy>(ctx, startNonce);
-		else
-		{
-			cryptonight_core_gpu_hash<CRYPTONIGHT_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight>(ctx, startNonce);
-		}
+		cryptonight_core_gpu_hash<CRYPTONIGHT_HEAVY_ITER, CRYPTONIGHT_HEAVY_MASK, CRYPTONIGHT_HEAVY_MEMORY/4, cryptonight_heavy>(ctx, startNonce);
 	}
 	else if(miner_algo == cryptonight)
 	{
@@ -504,10 +496,7 @@ void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t
 	}
 	else if(miner_algo == cryptonight_aeon)
 	{
-		if(version >= 7)
-			cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_aeon>(ctx, startNonce);
-		else
-			cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_lite>(ctx, startNonce);
+		cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_aeon>(ctx, startNonce);
 	}
 
 }
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
index c2a1f87..e2f0b2d 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
@@ -310,7 +310,7 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx)
 	return 1;
 }
 
-extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce, xmrstak_algo miner_algo, uint8_t version)
+extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce, xmrstak_algo miner_algo)
 {
 	int threadsperblock = 128;
 	uint32_t wsize = ctx->device_blocks * ctx->device_threads;
@@ -318,7 +318,7 @@ extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce
 	dim3 grid( ( wsize + threadsperblock - 1 ) / threadsperblock );
 	dim3 block( threadsperblock );
 
-	if(miner_algo == cryptonight_heavy && version >= 3)
+	if(miner_algo == cryptonight_heavy)
 	{
 		CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_extra_gpu_prepare<cryptonight_heavy><<<grid, block >>>( wsize, ctx->d_input, ctx->inputlen, startNonce,
 			ctx->d_ctx_state,ctx->d_ctx_state2, ctx->d_ctx_a, ctx->d_ctx_b, ctx->d_ctx_key1, ctx->d_ctx_key2 ));
@@ -333,7 +333,7 @@ extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce
 	}
 }
 
-extern "C" void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce,xmrstak_algo miner_algo, uint8_t version)
+extern "C" void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce,xmrstak_algo miner_algo)
 {
 	int threadsperblock = 128;
 	uint32_t wsize = ctx->device_blocks * ctx->device_threads;
@@ -344,7 +344,7 @@ extern "C" void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce,
 	CUDA_CHECK(ctx->device_id, cudaMemset( ctx->d_result_nonce, 0xFF, 10 * sizeof (uint32_t ) ));
 	CUDA_CHECK(ctx->device_id, cudaMemset( ctx->d_result_count, 0, sizeof (uint32_t ) ));
 
-	if(miner_algo == cryptonight_heavy && version >= 3)
+	if(miner_algo == cryptonight_heavy)
 	{
 		CUDA_CHECK_MSG_KERNEL(
 			ctx->device_id,
diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp
index bbaeb8b..5afc134 100644
--- a/xmrstak/jconf.cpp
+++ b/xmrstak/jconf.cpp
@@ -91,20 +91,21 @@ struct xmrstak_coin_algo
 	const char* coin_name;
 	xmrstak_algo algo;
 	xmrstak_algo algo_root;
+	uint8_t fork_height;
 	const char* default_pool;
 };
 
 xmrstak_coin_algo coin_algos[] = { 
-	{ "aeon7", cryptonight_aeon, cryptonight_lite, "mine.aeon-pool.com:5555" },
-	{ "cryptonight", cryptonight, cryptonight, nullptr },
-	{ "cryptonight_lite", cryptonight_lite, cryptonight_lite, nullptr },
-	{ "edollar", cryptonight, cryptonight, nullptr },
-	{ "electroneum", cryptonight, cryptonight, nullptr },
-	{ "graft", cryptonight, cryptonight, nullptr },
-	{ "intense", cryptonight, cryptonight, nullptr },
-	{ "karbo", cryptonight, cryptonight, nullptr },
-	{ "monero7", cryptonight_monero, cryptonight, "pool.usxmrpool.com:3333" },
-	{ "sumokoin", cryptonight_heavy, cryptonight, nullptr }
+	{ "aeon7", cryptonight_aeon, cryptonight_lite, 7u, "mine.aeon-pool.com:5555" },
+	{ "cryptonight", cryptonight, cryptonight, 0u, nullptr },
+	{ "cryptonight_lite", cryptonight_lite, cryptonight_lite, 0u, nullptr },
+	{ "edollar", cryptonight, cryptonight, 0u, nullptr },
+	{ "electroneum", cryptonight, cryptonight, 0u, nullptr },
+	{ "graft", cryptonight, cryptonight, 0u, nullptr },
+	{ "intense", cryptonight, cryptonight, 0u, nullptr },
+	{ "karbo", cryptonight, cryptonight, 0u, nullptr },
+	{ "monero7", cryptonight_monero, cryptonight, 7u, "pool.usxmrpool.com:3333" },
+	{ "sumokoin", cryptonight_heavy, cryptonight, 3u, nullptr }
 };
 
 constexpr size_t coin_alogo_size = (sizeof(coin_algos)/sizeof(coin_algos[0]));
@@ -632,6 +633,7 @@ bool jconf::parse_config(const char* sFilename, const char* sFilenamePools)
 		{
 			mining_algo = coin_algos[i].algo;
 			mining_algo_root = coin_algos[i].algo_root;
+			mining_fork_height = coin_algos[i].fork_height;
 			break;
 		}
 	}
diff --git a/xmrstak/jconf.hpp b/xmrstak/jconf.hpp
index f81a3a2..336e4be 100644
--- a/xmrstak/jconf.hpp
+++ b/xmrstak/jconf.hpp
@@ -48,9 +48,11 @@ public:
 
 	bool TlsSecureAlgos();
 
-	inline xmrstak_algo GetMiningAlgo() { return mining_algo; }
+	inline xmrstak_algo GetMiningAlgo() const { return mining_algo; }
 
-	inline xmrstak_algo GetMiningAlgoRoot() { return mining_algo_root; }
+	inline xmrstak_algo GetMiningAlgoRoot() const { return mining_algo_root; }
+
+	inline uint8_t GetMiningForkHeight() const { return mining_fork_height; }
 	
 	std::string GetMiningCoin();
 
@@ -94,4 +96,5 @@ private:
 	bool bHaveAes;
 	xmrstak_algo mining_algo;
 	xmrstak_algo mining_algo_root;
+	uint8_t mining_fork_height;
 };
-- 
GitLab