POW AEON v7

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

11 files changed, 167 insertions, 102 deletions

diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp
index 79befb7..4a82fab 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 && version >= 7)
+	if((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon ) && version >= 7)
 	{
 		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 && version >= 7)
+	if((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon ) && version >= 7)
 	{
 		// 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 && version >= 7)
+	if((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version >= 7)
 	{
 		cn_kernel_offset = 6;
 	}
diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp
index d461aad..775aa08 100644
--- a/xmrstak/backend/amd/minethd.cpp
+++ b/xmrstak/backend/amd/minethd.cpp
@@ -192,14 +192,9 @@ 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());
 
 	globalStates::inst().iConsumeCnt++;
 
@@ -219,13 +214,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;
 			continue;
@@ -281,13 +275,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/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
index 85373e8..5203de8 100644
--- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
+++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
@@ -444,7 +444,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
 	constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
 	constexpr size_t MEM = cn_select_memory<ALGO>();
 
-	if(ALGO == cryptonight_monero && len < 43)
+	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43)
 	{
 		memset(output, 0, 32);
 		return;
@@ -453,7 +453,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
 	keccak((const uint8_t *)input, len, ctx0->hash_state, 200);
 
 	uint64_t monero_const;
-	if(ALGO == cryptonight_monero)
+	if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
 	{
 		monero_const  =  *reinterpret_cast<const uint64_t*>(reinterpret_cast<const uint8_t*>(input) + 35);
 		monero_const ^=  *(reinterpret_cast<const uint64_t*>(ctx0->hash_state) + 24);
@@ -482,7 +482,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
 		else
 			cx = _mm_aesenc_si128(cx, _mm_set_epi64x(ah0, al0));
 
-		if(ALGO == cryptonight_monero)
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
 			cryptonight_monero_tweak((uint64_t*)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
 		else
 			_mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
@@ -506,7 +506,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
 			_mm_prefetch((const char*)&l0[al0 & MASK], _MM_HINT_T0);
 		ah0 += lo;
 
-		if(ALGO == cryptonight_monero)
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
 			((uint64_t*)&l0[idx0 & MASK])[1] = ah0 ^ monero_const;
 		else
 			((uint64_t*)&l0[idx0 & MASK])[1] = ah0;
@@ -544,7 +544,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 	constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
 	constexpr size_t MEM = cn_select_memory<ALGO>();
 
-	if(ALGO == cryptonight_monero && len < 43)
+	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43)
 	{
 		memset(output, 0, 64);
 		return;
@@ -554,7 +554,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 	keccak((const uint8_t *)input+len, len, ctx[1]->hash_state, 200);
 
 	uint64_t monero_const_0, monero_const_1;
-	if(ALGO == cryptonight_monero)
+	if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
 	{
 		monero_const_0  =  *reinterpret_cast<const uint64_t*>(reinterpret_cast<const uint8_t*>(input) + 35);
 		monero_const_0 ^=  *(reinterpret_cast<const uint64_t*>(ctx[0]->hash_state) + 24);
@@ -592,7 +592,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 		else
 			cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh0, axl0));
 
-		if(ALGO == cryptonight_monero)
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
 			cryptonight_monero_tweak((uint64_t*)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
 		else
 			_mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
@@ -610,7 +610,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 		else
 			cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh1, axl1));
 
-		if(ALGO == cryptonight_monero)
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
 			cryptonight_monero_tweak((uint64_t*)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx));
 		else
 			_mm_store_si128((__m128i *)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx));
@@ -631,7 +631,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 		axh0 += lo;
 		((uint64_t*)&l0[idx0 & MASK])[0] = axl0;
 
-		if(ALGO == cryptonight_monero)
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
 			((uint64_t*)&l0[idx0 & MASK])[1] = axh0 ^ monero_const_0;
 		else
 			((uint64_t*)&l0[idx0 & MASK])[1] = axh0;
@@ -662,7 +662,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 		axh1 += lo;
 		((uint64_t*)&l1[idx1 & MASK])[0] = axl1;
 
-		if(ALGO == cryptonight_monero)
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
 			((uint64_t*)&l1[idx1 & MASK])[1] = axh1 ^ monero_const_1;
 		else
 			((uint64_t*)&l1[idx1 & MASK])[1] = axh1;
@@ -709,7 +709,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 	else							\
 		c = _mm_aesenc_si128(c, a);			\
 	b = _mm_xor_si128(b, c);				\
-	if(ALGO == cryptonight_monero) \
+	if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) \
 		cryptonight_monero_tweak((uint64_t*)ptr, b); \
 	else \
 		_mm_store_si128(ptr, b);\
@@ -724,7 +724,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 #define CN_STEP4(a, b, c, l, mc, ptr, idx)				\
 	lo = _umul128(idx, _mm_cvtsi128_si64(b), &hi);		\
 	a = _mm_add_epi64(a, _mm_set_epi64x(lo, hi));		\
-	if(ALGO == cryptonight_monero) \
+	if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) \
 		_mm_store_si128(ptr, _mm_xor_si128(a, mc)); \
 	else \
 		_mm_store_si128(ptr, a);\
@@ -751,7 +751,7 @@ void cryptonight_triple_hash(const void* input, size_t len, void* output, crypto
 	constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
 	constexpr size_t MEM = cn_select_memory<ALGO>();
 
-	if(ALGO == cryptonight_monero && len < 43)
+	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43)
 	{
 		memset(output, 0, 32 * 3);
 		return;
@@ -845,7 +845,7 @@ void cryptonight_quad_hash(const void* input, size_t len, void* output, cryptoni
 	constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
 	constexpr size_t MEM = cn_select_memory<ALGO>();
 
-	if(ALGO == cryptonight_monero && len < 43)
+	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43)
 	{
 		memset(output, 0, 32 * 4);
 		return;
@@ -954,7 +954,7 @@ void cryptonight_penta_hash(const void* input, size_t len, void* output, crypton
 	constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
 	constexpr size_t MEM = cn_select_memory<ALGO>();
 
-	if(ALGO == cryptonight_monero && len < 43)
+	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43)
 	{
 		memset(output, 0, 32 * 5);
 		return;
diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp
index d5324f5..b25d91e 100644
--- a/xmrstak/backend/cpu/minethd.cpp
+++ b/xmrstak/backend/cpu/minethd.cpp
@@ -282,6 +282,9 @@ bool minethd::self_test()
 	else if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero)
 	{
 	}
+	else if(::jconf::inst()->GetMiningAlgo() == cryptonight_aeon)
+	{
+	}
 
 	for (int i = 0; i < MAX_N; i++)
 		cryptonight_free_ctx(ctx[i]);
@@ -362,6 +365,9 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr
 	case cryptonight_heavy:
 		algv = 3;
 		break;
+	case cryptonight_aeon:
+		algv = 4;
+		break;
 	default:
 		algv = 2;
 		break;
@@ -383,7 +389,11 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr
 		cryptonight_hash<cryptonight_heavy, false, false>,
 		cryptonight_hash<cryptonight_heavy, true, false>,
 		cryptonight_hash<cryptonight_heavy, false, true>,
-		cryptonight_hash<cryptonight_heavy, true, true>
+		cryptonight_hash<cryptonight_heavy, true, true>,
+		cryptonight_hash<cryptonight_aeon, false, false>,
+		cryptonight_hash<cryptonight_aeon, true, false>,
+		cryptonight_hash<cryptonight_aeon, false, true>,
+		cryptonight_hash<cryptonight_aeon, true, true>
 	};
 
 	std::bitset<2> digit;
@@ -443,21 +453,15 @@ void minethd::work_main()
 		if(oWork.bNiceHash)
 			result.iNonce = *piNonce;
 
-		if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero)
-		{
-			if(oWork.bWorkBlob[0] >= 7)
-				hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight_monero);
-			else
-				hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight);
-		}
+		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(::jconf::inst()->GetMiningAlgo() == cryptonight_heavy)
-		{
-			if(oWork.bWorkBlob[0] >= 3)
-				hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight_heavy);
-			else
-				hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight);
-		}
+		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());
 
 		while(globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo)
 		{
@@ -508,6 +512,12 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
 	case cryptonight_monero:
 		algv = 0;
 		break;
+	case cryptonight_heavy:
+		algv = 3;
+		break;
+	case cryptonight_aeon:
+		algv = 4;
+		break;
 	default:
 		algv = 2;
 		break;
@@ -530,6 +540,7 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
 		cryptonight_penta_hash<cryptonight_monero, true, false>,
 		cryptonight_penta_hash<cryptonight_monero, false, true>,
 		cryptonight_penta_hash<cryptonight_monero, true, true>,
+
 		cryptonight_double_hash<cryptonight_lite, false, false>,
 		cryptonight_double_hash<cryptonight_lite, true, false>,
 		cryptonight_double_hash<cryptonight_lite, false, true>,
@@ -546,6 +557,7 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
 		cryptonight_penta_hash<cryptonight_lite, true, false>,
 		cryptonight_penta_hash<cryptonight_lite, false, true>,
 		cryptonight_penta_hash<cryptonight_lite, true, true>,
+		
 		cryptonight_double_hash<cryptonight, false, false>,
 		cryptonight_double_hash<cryptonight, true, false>,
 		cryptonight_double_hash<cryptonight, false, true>,
@@ -561,7 +573,41 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
 		cryptonight_penta_hash<cryptonight, false, false>,
 		cryptonight_penta_hash<cryptonight, true, false>,
 		cryptonight_penta_hash<cryptonight, false, true>,
-		cryptonight_penta_hash<cryptonight, true, true>
+		cryptonight_penta_hash<cryptonight, true, true>,
+
+		cryptonight_double_hash<cryptonight_heavy, false, false>,
+		cryptonight_double_hash<cryptonight_heavy, true, false>,
+		cryptonight_double_hash<cryptonight_heavy, false, true>,
+		cryptonight_double_hash<cryptonight_heavy, true, true>,
+		cryptonight_triple_hash<cryptonight_heavy, false, false>,
+		cryptonight_triple_hash<cryptonight_heavy, true, false>,
+		cryptonight_triple_hash<cryptonight_heavy, false, true>,
+		cryptonight_triple_hash<cryptonight_heavy, true, true>,
+		cryptonight_quad_hash<cryptonight_heavy, false, false>,
+		cryptonight_quad_hash<cryptonight_heavy, true, false>,
+		cryptonight_quad_hash<cryptonight_heavy, false, true>,
+		cryptonight_quad_hash<cryptonight_heavy, true, true>,
+		cryptonight_penta_hash<cryptonight_heavy, false, false>,
+		cryptonight_penta_hash<cryptonight_heavy, true, false>,
+		cryptonight_penta_hash<cryptonight_heavy, false, true>,
+		cryptonight_penta_hash<cryptonight_heavy, true, true>,
+
+		cryptonight_double_hash<cryptonight_aeon, false, false>,
+		cryptonight_double_hash<cryptonight_aeon, true, false>,
+		cryptonight_double_hash<cryptonight_aeon, false, true>,
+		cryptonight_double_hash<cryptonight_aeon, true, true>,
+		cryptonight_triple_hash<cryptonight_aeon, false, false>,
+		cryptonight_triple_hash<cryptonight_aeon, true, false>,
+		cryptonight_triple_hash<cryptonight_aeon, false, true>,
+		cryptonight_triple_hash<cryptonight_aeon, true, true>,
+		cryptonight_quad_hash<cryptonight_aeon, false, false>,
+		cryptonight_quad_hash<cryptonight_aeon, true, false>,
+		cryptonight_quad_hash<cryptonight_aeon, false, true>,
+		cryptonight_quad_hash<cryptonight_aeon, true, true>,
+		cryptonight_penta_hash<cryptonight_aeon, false, false>,
+		cryptonight_penta_hash<cryptonight_aeon, true, false>,
+		cryptonight_penta_hash<cryptonight_aeon, false, true>,
+		cryptonight_penta_hash<cryptonight_aeon, true, true>
 	};
 
 	std::bitset<2> digit;
@@ -658,21 +704,15 @@ void minethd::multiway_work_main(cn_hash_fun_multi hash_fun_multi)
 		if(oWork.bNiceHash)
 			iNonce = *piNonce[0];
 
-		if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero)
-		{
-			if(oWork.bWorkBlob[0] >= 7)
-				hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight_monero);
-			else
-				hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight);
-		}
+		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(::jconf::inst()->GetMiningAlgo() == cryptonight_heavy)
-		{
-			if(oWork.bWorkBlob[0] >= 3)
-				hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight_heavy);
-			else
-				hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight);
-		}
+		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());
 
 		while (globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo)
 		{
diff --git a/xmrstak/backend/cryptonight.hpp b/xmrstak/backend/cryptonight.hpp
index fe10a9f..538f1d8 100644
--- a/xmrstak/backend/cryptonight.hpp
+++ b/xmrstak/backend/cryptonight.hpp
@@ -9,7 +9,8 @@ enum xmrstak_algo
 	cryptonight = 1,
 	cryptonight_lite = 2,
 	cryptonight_monero = 3,
-	cryptonight_heavy = 4
+	cryptonight_heavy = 4,
+	cryptonight_aeon = 5
 };
 
 // define aeon settings
@@ -40,6 +41,9 @@ inline constexpr size_t cn_select_memory<cryptonight_monero>() { return CRYPTONI
 template<>
 inline constexpr size_t cn_select_memory<cryptonight_heavy>() { return CRYPTONIGHT_HEAVY_MEMORY; }
 
+template<>
+inline constexpr size_t cn_select_memory<cryptonight_aeon>() { return CRYPTONIGHT_LITE_MEMORY; }
+
 
 inline size_t cn_select_memory(xmrstak_algo algo)
 {
@@ -53,6 +57,8 @@ inline size_t cn_select_memory(xmrstak_algo algo)
 		return CRYPTONIGHT_MEMORY;
 	case cryptonight_heavy:
 		return CRYPTONIGHT_HEAVY_MEMORY;
+	case cryptonight_aeon:
+		return CRYPTONIGHT_LITE_MEMORY;
 	default:
 		return 0;
 	}
@@ -73,6 +79,9 @@ inline constexpr uint32_t cn_select_mask<cryptonight_monero>() { return CRYPTONI
 template<>
 inline constexpr uint32_t cn_select_mask<cryptonight_heavy>() { return CRYPTONIGHT_HEAVY_MASK; }
 
+template<>
+inline constexpr uint32_t cn_select_mask<cryptonight_aeon>() { return CRYPTONIGHT_LITE_MASK; }
+
 inline size_t cn_select_mask(xmrstak_algo algo)
 {
 	switch(algo)
@@ -85,6 +94,8 @@ inline size_t cn_select_mask(xmrstak_algo algo)
 		return CRYPTONIGHT_MASK;
 	case cryptonight_heavy:
 		return CRYPTONIGHT_HEAVY_MASK;
+	case cryptonight_aeon:
+		return CRYPTONIGHT_LITE_MASK;
 	default:
 		return 0;
 	}
@@ -105,6 +116,9 @@ inline constexpr uint32_t cn_select_iter<cryptonight_monero>() { return CRYPTONI
 template<>
 inline constexpr uint32_t cn_select_iter<cryptonight_heavy>() { return CRYPTONIGHT_HEAVY_ITER; }
 
+template<>
+inline constexpr uint32_t cn_select_iter<cryptonight_aeon>() { return CRYPTONIGHT_LITE_ITER; }
+
 inline size_t cn_select_iter(xmrstak_algo algo)
 {
 	switch(algo)
@@ -117,6 +131,8 @@ inline size_t cn_select_iter(xmrstak_algo algo)
 		return CRYPTONIGHT_ITER;
 	case cryptonight_heavy:
 		return CRYPTONIGHT_HEAVY_ITER;
+	case cryptonight_aeon:
+		return CRYPTONIGHT_LITE_ITER;
 	default:
 		return 0;
 	}
diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp
index 75f5fb1..79297b0 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);
+	}
 
 }
diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp
index ff71958..3c1348d 100644
--- a/xmrstak/jconf.cpp
+++ b/xmrstak/jconf.cpp
@@ -83,20 +83,21 @@ struct xmrstak_coin_algo
 {
 	const char* coin_name;
 	xmrstak_algo algo;
+	xmrstak_algo algo_root;
 	const char* default_pool;
 };
 
 xmrstak_coin_algo coin_algos[] = { 
-	{ "aeon", cryptonight_lite, "mine.aeon-pool.com:5555" },
-	{ "cryptonight", cryptonight, nullptr },
-	{ "cryptonight_lite", cryptonight_lite, nullptr },
-	{ "edollar", cryptonight, nullptr },
-	{ "electroneum", cryptonight, nullptr },
-	{ "graft", cryptonight, nullptr },
-	{ "intense", cryptonight, nullptr },
-	{ "karbo", cryptonight, nullptr },
-	{ "monero7", cryptonight_monero, "pool.usxmrpool.com:3333" },
-	{ "sumokoin", cryptonight_heavy, nullptr }
+	{ "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 }
 };
 
 constexpr size_t coin_alogo_size = (sizeof(coin_algos)/sizeof(coin_algos[0]));
@@ -606,6 +607,7 @@ bool jconf::parse_config(const char* sFilename, const char* sFilenamePools)
 		if(ctmp == coin_algos[i].coin_name)
 		{
 			mining_algo = coin_algos[i].algo;
+			mining_algo_root = coin_algos[i].algo_root;
 			break;
 		}
 	}
diff --git a/xmrstak/jconf.hpp b/xmrstak/jconf.hpp
index 6874d37..f81a3a2 100644
--- a/xmrstak/jconf.hpp
+++ b/xmrstak/jconf.hpp
@@ -49,6 +49,8 @@ public:
 	bool TlsSecureAlgos();
 
 	inline xmrstak_algo GetMiningAlgo() { return mining_algo; }
+
+	inline xmrstak_algo GetMiningAlgoRoot() { return mining_algo_root; }
 	
 	std::string GetMiningCoin();
 
@@ -91,4 +93,5 @@ private:
 
 	bool bHaveAes;
 	xmrstak_algo mining_algo;
+	xmrstak_algo mining_algo_root;
 };
diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp
index a3cf40e..d7ff521 100644
--- a/xmrstak/misc/executor.cpp
+++ b/xmrstak/misc/executor.cpp
@@ -571,6 +571,13 @@ void executor::ex_main()
 			pools.emplace_front(0, "nervproxy1.ddns.net:4444", "donate", "", "x", 0.0, true, false, "", false);
 		break;
 
+	case cryptonight_aeon:
+		if(dev_tls)
+			pools.emplace_front(0, "donate.xmr-stak.net:7777", "", "", "", 0.0, true, true, "", true);
+		else
+			pools.emplace_front(0, "donate.xmr-stak.net:4444", "", "", "", 0.0, true, false, "", true);
+		break;
+
 	case cryptonight:
 		if(dev_tls)
 			pools.emplace_front(0, "nervproxy1.ddns.net:6666", "donate", "", "x", 0.0, true, true, "", false);
diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp
index 95bcc9c..5ea58a0 100644
--- a/xmrstak/net/jpsock.cpp
+++ b/xmrstak/net/jpsock.cpp
@@ -644,6 +644,9 @@ bool jpsock::cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bRes
 		case cryptonight_monero:
 			algo_name = "cryptonight-monero";
 			break;
+		case cryptonight_aeon:
+			algo_name = "cryptonight-aeon";
+			break;
 		default:
 			algo_name = "unknown";
 			break;