POW AEON v7

- add new pow for AEON
- fix missing cryptonight-heavy selection for multi hashes

2 files changed, 24 insertions, 23 deletions

diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp
index 153e4e3..f8e8409 100644
--- a/xmrstak/backend/nvidia/minethd.cpp
+++ b/xmrstak/backend/nvidia/minethd.cpp
@@ -238,14 +238,10 @@ void minethd::work_main()
 	cryptonight_ctx* cpu_ctx;
 	cpu_ctx = cpu::minethd::minethd_alloc_ctx();
 	auto miner_algo = ::jconf::inst()->GetMiningAlgo();
-	cn_hash_fun hash_fun;
-	if(miner_algo == cryptonight_monero || miner_algo == cryptonight_heavy)
-	{
-		// start with cryptonight and switch later if fork version is reached
-		hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight);
-	}
-	else
-		hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
+	
+	// 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());
+	
 	uint32_t iNonce;
 
 	globalStates::inst().iConsumeCnt++;
@@ -266,13 +262,12 @@ void minethd::work_main()
 
 			consume_work();
 			uint8_t new_version = oWork.getVersion();
-			if(miner_algo == cryptonight_monero && version < 7 && new_version >= 7)
+			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*/, cryptonight_monero);
-			}
-			else if(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3)
-			{
-				hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy);
+				hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
 			}
 			version = new_version;
 			continue;
@@ -335,13 +330,12 @@ void minethd::work_main()
 
 		consume_work();
 		uint8_t new_version = oWork.getVersion();
-		if(miner_algo == cryptonight_monero && version < 7 && new_version >= 7)
-		{
-			hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_monero);
-		}
-		else if(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3)
+		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*/, cryptonight_heavy);
+			hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
 		}
 		version = new_version;
 	}
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
index ede578f..a1a9e15 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
@@ -231,7 +231,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
 	uint32_t t1[2], t2[2], res;
 
 	uint32_t tweak1_2[2];
-	if (ALGO == cryptonight_monero)
+	if (ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
 	{
 		uint32_t * state = d_ctx_state + thread * 50;
 		tweak1_2[0] = (d_input[8] >> 24) | (d_input[9] << 8);
@@ -275,7 +275,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
 			t1[0] = shuffle<4>(sPtr,sub, d[x], 0);
 
 			const uint32_t z = d[0] ^ d[1];
-			if(ALGO == cryptonight_monero)
+			if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
 			{
 				const uint32_t table = 0x75310U;
 				const uint32_t index = ((z >> 26) & 12) | ((z >> 23) & 2);
@@ -304,7 +304,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
 			res = *( (uint64_t *) t2 )  >> ( sub & 1 ? 32 : 0 );
 
 			
-			if(ALGO == cryptonight_monero)
+			if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
 			{
 				const uint32_t tweaked_res = tweak1_2[sub & 1] ^ res;
 				const uint32_t long_state_update = sub2 ? tweaked_res : res;
@@ -502,5 +502,12 @@ void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t
 	{
 		cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_lite>(ctx, startNonce);
 	}
+	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);
+	}
 
 }