Merge pull request #67 from psychocrypt/topic-aeon2

add Aeon support

17 files changed, 307 insertions, 124 deletions

diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp
index 3575854..22ce5d0 100644
--- a/xmrstak/backend/amd/amd_gpu/gpu.cpp
+++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp
@@ -13,6 +13,9 @@
   * along with this program.  If not, see <http://www.gnu.org/licenses/>.
   */
 
+#include "xmrstak/backend/cryptonight.hpp"
+#include "xmrstak/jconf.hpp"
+
 #include <stdio.h>
 #include <string.h>
 #include <math.h>
@@ -245,8 +248,24 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_
 		return ERR_OCL_API;
 	}
 
+	size_t hashMemSize;
+	int threadMemMask;
+	int hasIterations;
+	if(::jconf::inst()->IsCurrencyMonero())
+	{
+		hashMemSize = MONERO_MEMORY;
+		threadMemMask = MONERO_MASK;
+		hasIterations = MONERO_ITER;
+	}
+	else
+	{
+		hashMemSize = AEON_MEMORY;
+		threadMemMask = AEON_MASK;
+		hasIterations = AEON_ITER;
+	}
+
 	size_t g_thd = ctx->rawIntensity;
-	ctx->ExtraBuffers[0] = clCreateBuffer(opencl_ctx, CL_MEM_READ_WRITE, (1 << 21) * g_thd, NULL, &ret);
+	ctx->ExtraBuffers[0] = clCreateBuffer(opencl_ctx, CL_MEM_READ_WRITE, hashMemSize * g_thd, NULL, &ret);
 	if(ret != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clCreateBuffer to create hash scratchpads buffer.", err_to_str(ret));
@@ -307,8 +326,9 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_
 		return ERR_OCL_API;
 	}
 
-	char options[32];
-	snprintf(options, sizeof(options), "-I. -DWORKSIZE=%llu", int_port(ctx->workSize));
+	char options[256];
+	snprintf(options, sizeof(options), 
+		"-DITERATIONS=%d -DMASK=%d -DWORKSIZE=%llu", hasIterations, threadMemMask, int_port(ctx->workSize));
 	ret = clBuildProgram(ctx->Program, 1, &ctx->DeviceID, options, NULL, NULL);
 	if(ret != CL_SUCCESS)
 	{
diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
index a6a5910..966199b 100644
--- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
+++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
@@ -91,7 +91,7 @@ static const __constant ulong keccakf_rndc[24] =
     0x8000000080008081, 0x8000000000008009, 0x000000000000008a,
     0x0000000000000088, 0x0000000080008009, 0x000000008000000a,
     0x000000008000808b, 0x800000000000008b, 0x8000000000008089,
-    0x8000000000008003, 0x8000000000008002, 0x8000000000000080, 
+    0x8000000000008003, 0x8000000000008002, 0x8000000000000080,
     0x000000000000800a, 0x800000008000000a, 0x8000000080008081,
     0x8000000000008080, 0x0000000080000001, 0x8000000080008008
 };
@@ -440,7 +440,7 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul
 	if(gIdx < Threads)
 	{
 		states += 25 * gIdx;
-		Scratchpad += gIdx * (0x80000 >> 2);
+		Scratchpad += gIdx * (ITERATIONS >> 2);
 
 		((ulong8 *)State)[0] = vload8(0, input);
 		State[8] = input[8];
@@ -482,7 +482,7 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul
 	if(gIdx < Threads)
 	{
 		#pragma unroll 2
-		for(int i = 0; i < 0x4000; ++i)
+		for(int i = 0; i < (ITERATIONS >> 5); ++i)
 		{
 			#pragma unroll
 			for(int j = 0; j < 10; ++j)
@@ -519,7 +519,7 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre
 	if(gIdx < Threads)
 	{
 		states += 25 * gIdx;
-		Scratchpad += gIdx * (0x80000 >> 2);
+		Scratchpad += gIdx * (ITERATIONS >> 2);
 
 		a[0] = states[0] ^ states[4];
 		b[0] = states[2] ^ states[6];
@@ -535,23 +535,23 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre
 	if(gIdx < Threads)
 	{
 		#pragma unroll 8
-		for(int i = 0; i < 0x80000; ++i)
+		for(int i = 0; i < ITERATIONS; ++i)
 		{
 			ulong c[2];
 
-			((uint4 *)c)[0] = Scratchpad[IDX((a[0] & 0x1FFFF0) >> 4)];
+			((uint4 *)c)[0] = Scratchpad[IDX((a[0] & MASK) >> 4)];
 			((uint4 *)c)[0] = AES_Round(AES0, AES1, AES2, AES3, ((uint4 *)c)[0], ((uint4 *)a)[0]);
 			//b_x ^= ((uint4 *)c)[0];
 
-			Scratchpad[IDX((a[0] & 0x1FFFF0) >> 4)] = b_x ^ ((uint4 *)c)[0];
+			Scratchpad[IDX((a[0] & MASK) >> 4)] = b_x ^ ((uint4 *)c)[0];
 
 			uint4 tmp;
-			tmp = Scratchpad[IDX((c[0] & 0x1FFFF0) >> 4)];
+			tmp = Scratchpad[IDX((c[0] & MASK) >> 4)];
 
 			a[1] += c[0] * as_ulong2(tmp).s0;
 			a[0] += mul_hi(c[0], as_ulong2(tmp).s0);
 
-			Scratchpad[IDX((c[0] & 0x1FFFF0) >> 4)] = ((uint4 *)a)[0];
+			Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0];
 
 			((uint4 *)a)[0] ^= tmp;
 
@@ -588,7 +588,7 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u
 	if(gIdx < Threads)
 	{
 		states += 25 * gIdx;
-		Scratchpad += gIdx * (0x80000 >> 2);
+		Scratchpad += gIdx * (ITERATIONS >> 2);
 
 		#if defined(__Tahiti__) || defined(__Pitcairn__)
 
@@ -611,7 +611,7 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u
 	if(gIdx < Threads)
 	{
 		#pragma unroll 2
-		for(int i = 0; i < 0x4000; ++i)
+		for(int i = 0; i < (ITERATIONS >> 5); ++i)
 		{
 			text ^= Scratchpad[IDX((i << 3) + get_local_id(1))];
 
diff --git a/xmrstak/backend/amd/autoAdjust.hpp b/xmrstak/backend/amd/autoAdjust.hpp
index 2a22a08..87e6299 100644
--- a/xmrstak/backend/amd/autoAdjust.hpp
+++ b/xmrstak/backend/amd/autoAdjust.hpp
@@ -8,6 +8,8 @@
 #include "xmrstak/misc/console.hpp"
 #include "xmrstak/misc/configEditor.hpp"
 #include "xmrstak/params.hpp"
+#include "xmrstak/backend/cryptonight.hpp"
+#include "xmrstak/jconf.hpp"
 
 #include <vector>
 #include <cstdio>
@@ -81,6 +83,16 @@ private:
 
 		constexpr size_t byteToMiB = 1024u * 1024u;
 
+		size_t hashMemSize;
+		if(::jconf::inst()->IsCurrencyMonero())
+		{
+			hashMemSize = MONERO_MEMORY;
+		}
+		else
+		{
+			hashMemSize = AEON_MEMORY;
+		}
+
 		std::string conf;
         int i = 0;
         for(auto& ctx : devVec)
@@ -88,7 +100,7 @@ private:
 			// keep 64MiB memory free (value is randomly chosen)
 			size_t availableMem = ctx.freeMem - (64u * 1024 * 1024);
 			// 224byte extra memory is used per thread for meta data
-			size_t perThread = (size_t(1u)<<21) + 224u;
+			size_t perThread = hashMemSize + 224u;
 			size_t max_intensity = availableMem / perThread;
 			// 1000 is a magic selected limit \todo select max intensity depending of the gpu type
 			size_t possibleIntensity = std::min( size_t(1000u) , max_intensity );
diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp
index e262e0e..f12e12c 100644
--- a/xmrstak/backend/amd/minethd.cpp
+++ b/xmrstak/backend/amd/minethd.cpp
@@ -183,7 +183,7 @@ void minethd::work_main()
 	uint64_t iCount = 0;
 	cryptonight_ctx* cpu_ctx;
 	cpu_ctx = cpu::minethd::minethd_alloc_ctx();
-	cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/);
+	cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, ::jconf::inst()->IsCurrencyMonero());
 	globalStates::inst().iConsumeCnt++;
 	
 	while (bQuit == 0)
diff --git a/xmrstak/backend/cpu/autoAdjust.hpp b/xmrstak/backend/cpu/autoAdjust.hpp
index 294bc6f..7bdb14e 100644
--- a/xmrstak/backend/cpu/autoAdjust.hpp
+++ b/xmrstak/backend/cpu/autoAdjust.hpp
@@ -6,6 +6,7 @@
 #include "xmrstak/jconf.hpp"
 #include "xmrstak/misc/configEditor.hpp"
 #include "xmrstak/params.hpp"
+#include "xmrstak/backend/cryptonight.hpp"
 #include <string>
 
 #ifdef _WIN32
@@ -32,8 +33,21 @@ class autoAdjust
 {
 public:
 
+	size_t hashMemSize;
+	size_t halfHashMemSize;
+
 	autoAdjust()
 	{
+		if(::jconf::inst()->IsCurrencyMonero())
+		{
+			hashMemSize = MONERO_MEMORY;
+			halfHashMemSize = hashMemSize / 2u;
+		}
+		else
+		{
+			hashMemSize = AEON_MEMORY;
+			halfHashMemSize = hashMemSize / 2u;
+		}
 	}
 
 	bool printConfig()
@@ -49,9 +63,9 @@ public:
 
 		std::string conf;
 
-		if(!detectL3Size() || L3KB_size < 1024 || L3KB_size > 102400)
+		if(!detectL3Size() || L3KB_size < halfHashMemSize || L3KB_size > (halfHashMemSize * 100u))
 		{
-			if(L3KB_size < 1024 || L3KB_size > 102400)
+			if(L3KB_size < halfHashMemSize || L3KB_size > (halfHashMemSize * 100))
 				printer::inst()->print_msg(L0, "Autoconf failed: L3 size sanity check failed - %u KB.", L3KB_size);
 
 			conf += std::string("    { \"low_power_mode\" : false, \"no_prefetch\" : true, \"affine_to_cpu\" : false },\n");
@@ -74,7 +88,7 @@ public:
 				if(L3KB_size <= 0)
 					break;
 
-				double_mode = L3KB_size / 2048 > (int32_t)(corecnt-i);
+				double_mode = L3KB_size / hashMemSize > (int32_t)(corecnt-i);
 
 				conf += std::string("    { \"low_power_mode\" : ");
 				conf += std::string(double_mode ? "true" : "false");
@@ -93,9 +107,9 @@ public:
 					aff_id++;
 
 				if(double_mode)
-					L3KB_size -= 4096;
+					L3KB_size -= hashMemSize * 2u;
 				else
-					L3KB_size -= 2048;
+					L3KB_size -= hashMemSize;
 			}
 		}
 
@@ -128,7 +142,7 @@ private:
 			}
 
 			L3KB_size = ((get_masked(cpu_info[1], 31, 22) + 1) * (get_masked(cpu_info[1], 21, 12) + 1) *
-				(get_masked(cpu_info[1], 11, 0) + 1) * (cpu_info[2] + 1)) / 1024;
+				(get_masked(cpu_info[1], 11, 0) + 1) * (cpu_info[2] + 1)) / halfHashMemSize;
 
 			return true;
 		}
diff --git a/xmrstak/backend/cpu/autoAdjustHwloc.hpp b/xmrstak/backend/cpu/autoAdjustHwloc.hpp
index ad3b863..ddeb89b 100644
--- a/xmrstak/backend/cpu/autoAdjustHwloc.hpp
+++ b/xmrstak/backend/cpu/autoAdjustHwloc.hpp
@@ -3,6 +3,7 @@
 #include "xmrstak/misc/console.hpp"
 #include "xmrstak/misc/configEditor.hpp"
 #include "xmrstak/params.hpp"
+#include "xmrstak/backend/cryptonight.hpp"
 
 #ifdef _WIN32
 #include <windows.h>
@@ -27,6 +28,16 @@ public:
 
 	autoAdjust()
 	{
+		if(::jconf::inst()->IsCurrencyMonero())
+		{
+			hashMemSize = MONERO_MEMORY;
+			halfHashMemSize = hashMemSize / 2u;
+		}
+		else
+		{
+			hashMemSize = AEON_MEMORY;
+			halfHashMemSize = hashMemSize / 2u;
+		}
 	}
 
 	bool printConfig()
@@ -86,7 +97,9 @@ public:
 	}
 
 private:
-	static constexpr size_t hashSize = 2 * 1024 * 1024;
+	size_t hashMemSize;
+	size_t halfHashMemSize;
+
 	std::vector<uint32_t> results;
 
 	template<typename func>
@@ -161,8 +174,8 @@ private:
 			{
 				hwloc_obj_t l2obj = obj->children[i];
 				//If L2 is exclusive and greater or equal to 2MB add room for one more hash
-				if(isCacheObject(l2obj) && l2obj->attr != nullptr && l2obj->attr->cache.size >= hashSize)
-					cacheSize += hashSize;
+				if(isCacheObject(l2obj) && l2obj->attr != nullptr && l2obj->attr->cache.size >= hashMemSize)
+					cacheSize += hashMemSize;
 			}
 		}
 
@@ -170,7 +183,7 @@ private:
 		cores.reserve(16);
 		findChildrenByType(obj, HWLOC_OBJ_CORE, [&cores](hwloc_obj_t found) { cores.emplace_back(found); } );
 
-		size_t cacheHashes = (cacheSize + hashSize/2) / hashSize;
+		size_t cacheHashes = (cacheSize + halfHashMemSize) / hashMemSize;
 
 		//Firstly allocate PU 0 of every CORE, then PU 1 etc.
 		size_t pu_id = 0;
diff --git a/xmrstak/backend/cpu/crypto/cryptonight.h b/xmrstak/backend/cpu/crypto/cryptonight.h
index 978c798..631c39a 100644
--- a/xmrstak/backend/cpu/crypto/cryptonight.h
+++ b/xmrstak/backend/cpu/crypto/cryptonight.h
@@ -7,8 +7,8 @@ extern "C" {
 
 #include <stddef.h>
 #include <inttypes.h>
+#include "xmrstak/backend/cryptonight.hpp"
 
-#define MEMORY  2097152
 
 typedef struct {
 	uint8_t hash_state[224]; // Need only 200, explicit align
diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
index 942d511..2a6a769 100644
--- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
+++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
@@ -287,7 +287,7 @@ void cn_implode_scratchpad(const __m128i* input, __m128i* output)
 	_mm_store_si128(output + 11, xout7);
 }
 
-template<size_t ITERATIONS, size_t MEM, bool SOFT_AES, bool PREFETCH>
+template<size_t MASK, size_t ITERATIONS, size_t MEM, bool SOFT_AES, bool PREFETCH>
 void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_ctx* ctx0)
 {
 	keccak((const uint8_t *)input, len, ctx0->hash_state, 200);
@@ -308,36 +308,36 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
 	for(size_t i = 0; i < ITERATIONS; i++)
 	{
 		__m128i cx;
-		cx = _mm_load_si128((__m128i *)&l0[idx0 & 0x1FFFF0]);
+		cx = _mm_load_si128((__m128i *)&l0[idx0 & MASK]);
 
 		if(SOFT_AES)
 			cx = soft_aesenc(cx, _mm_set_epi64x(ah0, al0));
 		else
 			cx = _mm_aesenc_si128(cx, _mm_set_epi64x(ah0, al0));
 
-		_mm_store_si128((__m128i *)&l0[idx0 & 0x1FFFF0], _mm_xor_si128(bx0, cx));
+		_mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
 		idx0 = _mm_cvtsi128_si64(cx);
 		bx0 = cx;
 
 		if(PREFETCH)
-			_mm_prefetch((const char*)&l0[idx0 & 0x1FFFF0], _MM_HINT_T0);
+			_mm_prefetch((const char*)&l0[idx0 & MASK], _MM_HINT_T0);
 
 		uint64_t hi, lo, cl, ch;
-		cl = ((uint64_t*)&l0[idx0 & 0x1FFFF0])[0];
-		ch = ((uint64_t*)&l0[idx0 & 0x1FFFF0])[1];
+		cl = ((uint64_t*)&l0[idx0 & MASK])[0];
+		ch = ((uint64_t*)&l0[idx0 & MASK])[1];
 
 		lo = _umul128(idx0, cl, &hi);
 
 		al0 += hi;
 		ah0 += lo;
-		((uint64_t*)&l0[idx0 & 0x1FFFF0])[0] = al0;
-		((uint64_t*)&l0[idx0 & 0x1FFFF0])[1] = ah0;
+		((uint64_t*)&l0[idx0 & MASK])[0] = al0;
+		((uint64_t*)&l0[idx0 & MASK])[1] = ah0;
 		ah0 ^= ch;
 		al0 ^= cl;
 		idx0 = al0;
 
 		if(PREFETCH)
-			_mm_prefetch((const char*)&l0[idx0 & 0x1FFFF0], _MM_HINT_T0);
+			_mm_prefetch((const char*)&l0[idx0 & MASK], _MM_HINT_T0);
 	}
 
 	// Optim - 90% time boundary
@@ -352,7 +352,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
 // This lovely creation will do 2 cn hashes at a time. We have plenty of space on silicon
 // to fit temporary vars for two contexts. Function will read len*2 from input and write 64 bytes to output
 // We are still limited by L3 cache, so doubling will only work with CPUs where we have more than 2MB to core (Xeons)
-template<size_t ITERATIONS, size_t MEM, bool SOFT_AES, bool PREFETCH>
+template<size_t MASK, size_t ITERATIONS, size_t MEM, bool SOFT_AES, bool PREFETCH>
 void cryptonight_double_hash(const void* input, size_t len, void* output, cryptonight_ctx* __restrict ctx0, cryptonight_ctx* __restrict ctx1)
 {
 	keccak((const uint8_t *)input, len, ctx0->hash_state, 200);
@@ -381,66 +381,66 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
 	for (size_t i = 0; i < ITERATIONS; i++)
 	{
 		__m128i cx;
-		cx = _mm_load_si128((__m128i *)&l0[idx0 & 0x1FFFF0]);
+		cx = _mm_load_si128((__m128i *)&l0[idx0 & MASK]);
 
 		if(SOFT_AES)
 			cx = soft_aesenc(cx, _mm_set_epi64x(axh0, axl0));
 		else
 			cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh0, axl0));
 
-		_mm_store_si128((__m128i *)&l0[idx0 & 0x1FFFF0], _mm_xor_si128(bx0, cx));
+		_mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
 		idx0 = _mm_cvtsi128_si64(cx);
 		bx0 = cx;
 
 		if(PREFETCH)
-			_mm_prefetch((const char*)&l0[idx0 & 0x1FFFF0], _MM_HINT_T0);
+			_mm_prefetch((const char*)&l0[idx0 & MASK], _MM_HINT_T0);
 
-		cx = _mm_load_si128((__m128i *)&l1[idx1 & 0x1FFFF0]);
+		cx = _mm_load_si128((__m128i *)&l1[idx1 & MASK]);
 
 		if(SOFT_AES)
 			cx = soft_aesenc(cx, _mm_set_epi64x(axh1, axl1));
 		else
 			cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh1, axl1));
 
-		_mm_store_si128((__m128i *)&l1[idx1 & 0x1FFFF0], _mm_xor_si128(bx1, cx));
+		_mm_store_si128((__m128i *)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx));
 		idx1 = _mm_cvtsi128_si64(cx);
 		bx1 = cx;
 
 		if(PREFETCH)
-			_mm_prefetch((const char*)&l1[idx1 & 0x1FFFF0], _MM_HINT_T0);
+			_mm_prefetch((const char*)&l1[idx1 & MASK], _MM_HINT_T0);
 
 		uint64_t hi, lo, cl, ch;
-		cl = ((uint64_t*)&l0[idx0 & 0x1FFFF0])[0];
-		ch = ((uint64_t*)&l0[idx0 & 0x1FFFF0])[1];
+		cl = ((uint64_t*)&l0[idx0 & MASK])[0];
+		ch = ((uint64_t*)&l0[idx0 & MASK])[1];
 
 		lo = _umul128(idx0, cl, &hi);
 
 		axl0 += hi;
 		axh0 += lo;
-		((uint64_t*)&l0[idx0 & 0x1FFFF0])[0] = axl0;
-		((uint64_t*)&l0[idx0 & 0x1FFFF0])[1] = axh0;
+		((uint64_t*)&l0[idx0 & MASK])[0] = axl0;
+		((uint64_t*)&l0[idx0 & MASK])[1] = axh0;
 		axh0 ^= ch;
 		axl0 ^= cl;
 		idx0 = axl0;
 
 		if(PREFETCH)
-			_mm_prefetch((const char*)&l0[idx0 & 0x1FFFF0], _MM_HINT_T0);
+			_mm_prefetch((const char*)&l0[idx0 & MASK], _MM_HINT_T0);
 
-		cl = ((uint64_t*)&l1[idx1 & 0x1FFFF0])[0];
-		ch = ((uint64_t*)&l1[idx1 & 0x1FFFF0])[1];
+		cl = ((uint64_t*)&l1[idx1 & MASK])[0];
+		ch = ((uint64_t*)&l1[idx1 & MASK])[1];
 
 		lo = _umul128(idx1, cl, &hi);
 
 		axl1 += hi;
 		axh1 += lo;
-		((uint64_t*)&l1[idx1 & 0x1FFFF0])[0] = axl1;
-		((uint64_t*)&l1[idx1 & 0x1FFFF0])[1] = axh1;
+		((uint64_t*)&l1[idx1 & MASK])[0] = axl1;
+		((uint64_t*)&l1[idx1 & MASK])[1] = axh1;
 		axh1 ^= ch;
 		axl1 ^= cl;
 		idx1 = axl1;
 
 		if(PREFETCH)
-			_mm_prefetch((const char*)&l1[idx1 & 0x1FFFF0], _MM_HINT_T0);
+			_mm_prefetch((const char*)&l1[idx1 & MASK], _MM_HINT_T0);
 	}
 
 	// Optim - 90% time boundary
diff --git a/xmrstak/backend/cpu/crypto/cryptonight_common.cpp b/xmrstak/backend/cpu/crypto/cryptonight_common.cpp
index 0690415..c73dbd8 100644
--- a/xmrstak/backend/cpu/crypto/cryptonight_common.cpp
+++ b/xmrstak/backend/cpu/crypto/cryptonight_common.cpp
@@ -30,6 +30,8 @@ extern "C"
 }
 #include "cryptonight.h"
 #include "cryptonight_aesni.h"
+#include "xmrstak/backend/cryptonight.hpp"
+#include "xmrstak/jconf.hpp"
 #include <stdio.h>
 #include <stdlib.h>
 
@@ -194,12 +196,21 @@ size_t cryptonight_init(size_t use_fast_mem, size_t use_mlock, alloc_msg* msg)
 
 cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, alloc_msg* msg)
 {
+	size_t hashMemSize;
+	if(::jconf::inst()->IsCurrencyMonero())
+	{
+		hashMemSize = MONERO_MEMORY;
+	}
+	else
+	{
+		hashMemSize = AEON_MEMORY;
+	}
 	cryptonight_ctx* ptr = (cryptonight_ctx*)_mm_malloc(sizeof(cryptonight_ctx), 4096);
 
 	if(use_fast_mem == 0)
 	{
 		// use 2MiB aligned memory
-		ptr->long_state = (uint8_t*)_mm_malloc(MEMORY, 2*1024*1024);
+		ptr->long_state = (uint8_t*)_mm_malloc(hashMemSize, hashMemSize);
 		ptr->ctx_info[0] = 0;
 		ptr->ctx_info[1] = 0;
 		return ptr;
@@ -208,7 +219,7 @@ cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, al
 #ifdef _WIN32
 	SIZE_T iLargePageMin = GetLargePageMinimum();
 
-	if(MEMORY > iLargePageMin)
+	if(hashMemSize > iLargePageMin)
 		iLargePageMin *= 2;
 
 	ptr->long_state = (uint8_t*)VirtualAlloc(NULL, iLargePageMin,
@@ -231,13 +242,13 @@ cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, al
 #else
 
 #if defined(__APPLE__)
-	ptr->long_state  = (uint8_t*)mmap(0, MEMORY, PROT_READ | PROT_WRITE,
+	ptr->long_state  = (uint8_t*)mmap(0, hashMemSize, PROT_READ | PROT_WRITE,
 		MAP_PRIVATE | MAP_ANON, VM_FLAGS_SUPERPAGE_SIZE_2MB, 0);
 #elif defined(__FreeBSD__)
-	ptr->long_state = (uint8_t*)mmap(0, MEMORY, PROT_READ | PROT_WRITE,
+	ptr->long_state = (uint8_t*)mmap(0, hashMemSize, PROT_READ | PROT_WRITE,
 		MAP_PRIVATE | MAP_ANONYMOUS | MAP_ALIGNED_SUPER | MAP_PREFAULT_READ, -1, 0);
 #else
-	ptr->long_state = (uint8_t*)mmap(0, MEMORY, PROT_READ | PROT_WRITE,
+	ptr->long_state = (uint8_t*)mmap(0, hashMemSize, PROT_READ | PROT_WRITE,
 		MAP_PRIVATE | MAP_ANONYMOUS | MAP_HUGETLB | MAP_POPULATE, 0, 0);
 #endif
 
@@ -250,11 +261,11 @@ cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, al
 
 	ptr->ctx_info[0] = 1;
 
-	if(madvise(ptr->long_state, MEMORY, MADV_RANDOM|MADV_WILLNEED) != 0)
+	if(madvise(ptr->long_state, hashMemSize, MADV_RANDOM|MADV_WILLNEED) != 0)
 		msg->warning = "madvise failed";
 
 	ptr->ctx_info[1] = 0;
-	if(use_mlock != 0 && mlock(ptr->long_state, MEMORY) != 0)
+	if(use_mlock != 0 && mlock(ptr->long_state, hashMemSize) != 0)
 		msg->warning = "mlock failed";
 	else
 		ptr->ctx_info[1] = 1;
@@ -265,14 +276,23 @@ cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, al
 
 void cryptonight_free_ctx(cryptonight_ctx* ctx)
 {
+	size_t hashMemSize;
+	if(::jconf::inst()->IsCurrencyMonero())
+	{
+		hashMemSize = MONERO_MEMORY;
+	}
+	else
+	{
+		hashMemSize = AEON_MEMORY;
+	}
 	if(ctx->ctx_info[0] != 0)
 	{
 #ifdef _WIN32
 		VirtualFree(ctx->long_state, 0, MEM_RELEASE);
 #else
 		if(ctx->ctx_info[1] != 0)
-			munlock(ctx->long_state, MEMORY);
-		munmap(ctx->long_state, MEMORY);
+			munlock(ctx->long_state, hashMemSize);
+		munmap(ctx->long_state, hashMemSize);
 #endif // _WIN32
 	}
 	else
diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp
index 1524eca..d36ebf1 100644
--- a/xmrstak/backend/cpu/minethd.cpp
+++ b/xmrstak/backend/cpu/minethd.cpp
@@ -199,30 +199,34 @@ bool minethd::self_test()
 		return false;
 	}
 
-	unsigned char out[64];
-	bool bResult;
+	bool bResult = true;
 
-	cn_hash_fun hashf;
-	cn_hash_fun_dbl hashdf;
+	bool mineMonero = ::jconf::inst()->IsCurrencyMonero();
+	if(mineMonero)
+	{
+		unsigned char out[64];
+		cn_hash_fun hashf;
+		cn_hash_fun_dbl hashdf;
 
-	hashf = func_selector(::jconf::inst()->HaveHardwareAes(), false);
-	hashf("This is a test", 14, out, ctx0);
-	bResult = memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 32) == 0;
 
-	hashf = func_selector(::jconf::inst()->HaveHardwareAes(), true);
-	hashf("This is a test", 14, out, ctx0);
-	bResult &= memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 32) == 0;
+		hashf = func_selector(::jconf::inst()->HaveHardwareAes(), false, mineMonero);
+		hashf("This is a test", 14, out, ctx0);
+		bResult = memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 32) == 0;
 
-	hashdf = func_dbl_selector(::jconf::inst()->HaveHardwareAes(), false);
-	hashdf("The quick brown fox jumps over the lazy dogThe quick brown fox jumps over the lazy log", 43, out, ctx0, ctx1);
-	bResult &= memcmp(out, "\x3e\xbb\x7f\x9f\x7d\x27\x3d\x7c\x31\x8d\x86\x94\x77\x55\x0c\xc8\x00\xcf\xb1\x1b\x0c\xad\xb7\xff\xbd\xf6\xf8\x9f\x3a\x47\x1c\x59"
-		                   "\xb4\x77\xd5\x02\xe4\xd8\x48\x7f\x42\xdf\xe3\x8e\xed\x73\x81\x7a\xda\x91\xb7\xe2\x63\xd2\x91\x71\xb6\x5c\x44\x3a\x01\x2a\x41\x22", 64) == 0;
+		hashf = func_selector(::jconf::inst()->HaveHardwareAes(), true, mineMonero);
+		hashf("This is a test", 14, out, ctx0);
+		bResult &= memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 32) == 0;
 
-	hashdf = func_dbl_selector(::jconf::inst()->HaveHardwareAes(), true);
-	hashdf("The quick brown fox jumps over the lazy dogThe quick brown fox jumps over the lazy log", 43, out, ctx0, ctx1);
-	bResult &= memcmp(out, "\x3e\xbb\x7f\x9f\x7d\x27\x3d\x7c\x31\x8d\x86\x94\x77\x55\x0c\xc8\x00\xcf\xb1\x1b\x0c\xad\xb7\xff\xbd\xf6\xf8\x9f\x3a\x47\x1c\x59"
-		                   "\xb4\x77\xd5\x02\xe4\xd8\x48\x7f\x42\xdf\xe3\x8e\xed\x73\x81\x7a\xda\x91\xb7\xe2\x63\xd2\x91\x71\xb6\x5c\x44\x3a\x01\x2a\x41\x22", 64) == 0;
+		hashdf = func_dbl_selector(::jconf::inst()->HaveHardwareAes(), false, mineMonero);
+		hashdf("The quick brown fox jumps over the lazy dogThe quick brown fox jumps over the lazy log", 43, out, ctx0, ctx1);
+		bResult &= memcmp(out, "\x3e\xbb\x7f\x9f\x7d\x27\x3d\x7c\x31\x8d\x86\x94\x77\x55\x0c\xc8\x00\xcf\xb1\x1b\x0c\xad\xb7\xff\xbd\xf6\xf8\x9f\x3a\x47\x1c\x59"
+							   "\xb4\x77\xd5\x02\xe4\xd8\x48\x7f\x42\xdf\xe3\x8e\xed\x73\x81\x7a\xda\x91\xb7\xe2\x63\xd2\x91\x71\xb6\x5c\x44\x3a\x01\x2a\x41\x22", 64) == 0;
 
+		hashdf = func_dbl_selector(::jconf::inst()->HaveHardwareAes(), true, mineMonero);
+		hashdf("The quick brown fox jumps over the lazy dogThe quick brown fox jumps over the lazy log", 43, out, ctx0, ctx1);
+		bResult &= memcmp(out, "\x3e\xbb\x7f\x9f\x7d\x27\x3d\x7c\x31\x8d\x86\x94\x77\x55\x0c\xc8\x00\xcf\xb1\x1b\x0c\xad\xb7\xff\xbd\xf6\xf8\x9f\x3a\x47\x1c\x59"
+							   "\xb4\x77\xd5\x02\xe4\xd8\x48\x7f\x42\xdf\xe3\x8e\xed\x73\x81\x7a\xda\x91\xb7\xe2\x63\xd2\x91\x71\xb6\x5c\x44\x3a\x01\x2a\x41\x22", 64) == 0;
+	}
 	cryptonight_free_ctx(ctx0);
 	cryptonight_free_ctx(ctx1);
 
@@ -285,24 +289,47 @@ void minethd::consume_work()
 	globalStates::inst().inst().iConsumeCnt++;
 }
 
-minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch)
+minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, bool mineMonero)
 {
 	// We have two independent flag bits in the functions
 	// therefore we will build a binary digit and select the
 	// function as a two digit binary
-	// Digit order SOFT_AES, NO_PREFETCH
-
-	static const cn_hash_fun func_table[4] = {
-		cryptonight_hash<0x80000, MEMORY, false, false>,
-		cryptonight_hash<0x80000, MEMORY, false, true>,
-		cryptonight_hash<0x80000, MEMORY, true, false>,
-		cryptonight_hash<0x80000, MEMORY, true, true>
+	// Digit order SOFT_AES, NO_PREFETCH, MINER_ALGO
+
+	static const cn_hash_fun func_table[] = {
+		/* there will be 8 function entries if `CONF_NO_MONERO` and `CONF_NO_AEON`
+		 * is not defined. If one is defined there will be 4 entries.
+		 */
+#ifndef CONF_NO_MONERO
+		cryptonight_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, false>,
+		cryptonight_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, true>,
+		cryptonight_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, false>,
+		cryptonight_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, true>
+#endif
+#if (!defined(CONF_NO_AEON)) && (!defined(CONF_NO_MONERO))
+		// comma will be added only if Monero and Aeon is build
+		,
+#endif
+#ifndef CONF_NO_AEON
+		cryptonight_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, false>,
+		cryptonight_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, true>,
+		cryptonight_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, false>,
+		cryptonight_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, true>
+#endif
 	};
 
-	std::bitset<2> digit;
+	std::bitset<3> digit;
 	digit.set(0, !bNoPrefetch);
 	digit.set(1, !bHaveAes);
 
+	// define aeon settings
+#if defined(CONF_NO_AEON) || defined(CONF_NO_MONERO)
+	// ignore 3rd bit if only on currency is active
+	digit.set(2, 0);
+#else
+	digit.set(2, !mineMonero);
+#endif
+
 	return func_table[digit.to_ulong()];
 }
 
@@ -320,7 +347,7 @@ void minethd::work_main()
 	uint32_t* piNonce;
 	job_result result;
 
-	hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch);
+	hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->IsCurrencyMonero());
 	ctx = minethd_alloc_ctx();
 
 	piHashVal = (uint64_t*)(result.bResult + 24);
@@ -382,24 +409,47 @@ void minethd::work_main()
 	cryptonight_free_ctx(ctx);
 }
 
-minethd::cn_hash_fun_dbl minethd::func_dbl_selector(bool bHaveAes, bool bNoPrefetch)
+minethd::cn_hash_fun_dbl minethd::func_dbl_selector(bool bHaveAes, bool bNoPrefetch, bool mineMonero)
 {
 	// We have two independent flag bits in the functions
 	// therefore we will build a binary digit and select the
 	// function as a two digit binary
-	// Digit order SOFT_AES, NO_PREFETCH
-
-	static const cn_hash_fun_dbl func_table[4] = {
-		cryptonight_double_hash<0x80000, MEMORY, false, false>,
-		cryptonight_double_hash<0x80000, MEMORY, false, true>,
-		cryptonight_double_hash<0x80000, MEMORY, true, false>,
-		cryptonight_double_hash<0x80000, MEMORY, true, true>
+	// Digit order SOFT_AES, NO_PREFETCH, MINER_ALGO
+
+	static const cn_hash_fun_dbl func_table[] = {
+		/* there will be 8 function entries if `CONF_NO_MONERO` and `CONF_NO_AEON`
+		 * is not defined. If one is defined there will be 4 entries.
+		 */
+#ifndef CONF_NO_MONERO
+		cryptonight_double_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, false>,
+		cryptonight_double_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, true>,
+		cryptonight_double_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, false>,
+		cryptonight_double_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, true>
+#endif
+#if (!defined(CONF_NO_AEON)) && (!defined(CONF_NO_MONERO))
+		// comma will be added only if Monero and Aeon is build
+		,
+#endif
+#ifndef CONF_NO_AEON
+		cryptonight_double_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, false>,
+		cryptonight_double_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, true>,
+		cryptonight_double_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, false>,
+		cryptonight_double_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, true>
+#endif
 	};
 
-	std::bitset<2> digit;
+	std::bitset<3> digit;
 	digit.set(0, !bNoPrefetch);
 	digit.set(1, !bHaveAes);
 
+	// define aeon settings
+#if defined(CONF_NO_AEON) || defined(CONF_NO_MONERO)
+	// ignore 3rd bit if only on currency is active
+	digit.set(2, 0);
+#else
+	digit.set(2, !mineMonero);
+#endif
+
 	return func_table[digit.to_ulong()];
 }
 
@@ -428,7 +478,7 @@ void minethd::double_work_main()
 	uint32_t iNonce;
 	job_result res;
 
-	hash_fun = func_dbl_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch);
+	hash_fun = func_dbl_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->IsCurrencyMonero());
 	ctx0 = minethd_alloc_ctx();
 	ctx1 = minethd_alloc_ctx();
 
diff --git a/xmrstak/backend/cpu/minethd.hpp b/xmrstak/backend/cpu/minethd.hpp
index 5ffd44e..670ec8d 100644
--- a/xmrstak/backend/cpu/minethd.hpp
+++ b/xmrstak/backend/cpu/minethd.hpp
@@ -23,7 +23,7 @@ public:
 
 	typedef void (*cn_hash_fun)(const void*, size_t, void*, cryptonight_ctx*);
 
-	static cn_hash_fun func_selector(bool bHaveAes, bool bNoPrefetch);
+	static cn_hash_fun func_selector(bool bHaveAes, bool bNoPrefetch, bool mineMonero);
 	static bool thd_setaffinity(std::thread::native_handle_type h, uint64_t cpu_id);
 
 	static cryptonight_ctx* minethd_alloc_ctx();
@@ -31,7 +31,7 @@ public:
 private:
 
 	typedef void (*cn_hash_fun_dbl)(const void*, size_t, void*, cryptonight_ctx* __restrict, cryptonight_ctx* __restrict);
-	static cn_hash_fun_dbl func_dbl_selector(bool bHaveAes, bool bNoPrefetch);
+	static cn_hash_fun_dbl func_dbl_selector(bool bHaveAes, bool bNoPrefetch, bool mineMonero);
 
 	minethd(miner_work& pWork, size_t iNo, bool double_work, bool no_prefetch, int64_t affinity);
 
diff --git a/xmrstak/backend/cryptonight.hpp b/xmrstak/backend/cryptonight.hpp
new file mode 100644
index 0000000..0ef5ae7
--- /dev/null
+++ b/xmrstak/backend/cryptonight.hpp
@@ -0,0 +1,12 @@
+#pragma once
+
+// define aeon settings
+#define AEON_MEMORY 1048576llu
+#define AEON_MASK 0xFFFF0
+#define AEON_ITER 0x40000
+
+// define xmr settings
+#define MONERO_MEMORY 2097152llu
+#define MONERO_MASK 0x1FFFF0
+#define MONERO_ITER 0x80000
+
diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp
index cf58a21..0bc6214 100644
--- a/xmrstak/backend/nvidia/minethd.cpp
+++ b/xmrstak/backend/nvidia/minethd.cpp
@@ -32,6 +32,8 @@
 #include "xmrstak/jconf.hpp"
 #include "xmrstak/misc/environment.hpp"
 #include "xmrstak/backend/cpu/hwlocMemory.hpp"
+#include "xmrstak/backend/cryptonight.hpp"
+#include "xmrstak/misc/utility.hpp"
 
 #include <assert.h>
 #include <cmath>
@@ -208,7 +210,7 @@ void minethd::work_main()
 	uint64_t iCount = 0;
 	cryptonight_ctx* cpu_ctx;
 	cpu_ctx = cpu::minethd::minethd_alloc_ctx();
-	cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/);
+	cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, ::jconf::inst()->IsCurrencyMonero());
 	uint32_t iNonce;
 
 	globalStates::inst().iConsumeCnt++;
@@ -218,6 +220,9 @@ void minethd::work_main()
 		printer::inst()->print_msg(L0, "Setup failed for GPU %d. Exitting.\n", (int)iThreadNo);
 		std::exit(0);
 	}
+
+	bool mineMonero = strcmp_i(::jconf::inst()->GetCurrency(), "monero");
+	bool useAEON = strcmp_i(::jconf::inst()->GetCurrency(), "aeon");
 	
 	while (bQuit == 0)
 	{
@@ -256,7 +261,18 @@ void minethd::work_main()
 			uint32_t foundCount;
 
 			cryptonight_extra_cpu_prepare(&ctx, iNonce);
-			cryptonight_core_cpu_hash(&ctx);
+#ifndef CONF_NO_MONERO
+			if(mineMonero)
+			{
+				cryptonight_core_cpu_hash<MONERO_ITER, MONERO_MASK, 19>(&ctx);
+			}
+#endif
+#ifndef CONF_NO_AEON
+			if(useAEON)
+			{
+				cryptonight_core_cpu_hash<MONERO_ITER, MONERO_MASK, 18>(&ctx);
+			}
+#endif
 			cryptonight_extra_cpu_final(&ctx, iNonce, oWork.iTarget, &foundCount, foundNonce);
 
 			for(size_t i = 0; i < foundCount; i++)
diff --git a/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp b/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp
index 784c38d..4e0ace7 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp
+++ b/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp
@@ -41,8 +41,9 @@ 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);
-void cryptonight_core_cpu_hash(nvid_ctx* ctx);
 void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce);
-
 }
 
+template<size_t ITERATIONS, size_t THREAD_SHIFT, uint32_t MASK>
+void cryptonight_core_cpu_hash(nvid_ctx* ctx);
+
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
index e1b78ce..5bea230 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
@@ -92,6 +92,7 @@ __device__ __forceinline__ void storeGlobal32( T* addr, T const & val )
 	asm volatile( "st.global.cg.u32 [%0], %1;" : : "l"( addr ), "r"( val ) );
 }
 
+template<size_t ITERATIONS, size_t THREAD_SHIFT>
 __global__ void cryptonight_core_gpu_phase1( int threads, int bfactor, int partidx, uint32_t * __restrict__ long_state, uint32_t * __restrict__ ctx_state, uint32_t * __restrict__ ctx_key1 )
 {
 	__shared__ uint32_t sharedMemory[1024];
@@ -102,7 +103,7 @@ __global__ void cryptonight_core_gpu_phase1( int threads, int bfactor, int parti
 	const int thread = ( blockDim.x * blockIdx.x + threadIdx.x ) >> 3;
 	const int sub = ( threadIdx.x & 7 ) << 2;
 
-	const int batchsize = 0x80000 >> bfactor;
+	const int batchsize = ITERATIONS >> bfactor;
 	const int start = partidx * batchsize;
 	const int end = start + batchsize;
 
@@ -121,13 +122,13 @@ __global__ void cryptonight_core_gpu_phase1( int threads, int bfactor, int parti
 	else
 	{
 		// load previous text data
-		MEMCPY8( text, &long_state[( (uint64_t) thread << 19 ) + sub + start - 32], 2 );
+		MEMCPY8( text, &long_state[( (uint64_t) thread << THREAD_SHIFT ) + sub + start - 32], 2 );
 	}
 	__syncthreads( );
 	for ( int i = start; i < end; i += 32 )
 	{
 		cn_aes_pseudo_round_mut( sharedMemory, text, key );
-		MEMCPY8(&long_state[((uint64_t) thread << 19) + (sub + i)], text, 2);
+		MEMCPY8(&long_state[((uint64_t) thread << THREAD_SHIFT) + (sub + i)], text, 2);
 	}
 }
 
@@ -167,6 +168,7 @@ __forceinline__ __device__ uint32_t shuffle(volatile uint32_t* ptr,const uint32_
 #ifdef XMR_STAK_THREADS
 __launch_bounds__( XMR_STAK_THREADS * 4 )
 #endif
+template<size_t ITERATIONS, size_t THREAD_SHIFT, uint32_t MASK>
 __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int partidx, uint32_t * d_long_state, uint32_t * d_ctx_a, uint32_t * d_ctx_b )
 {
 	__shared__ uint32_t sharedMemory[1024];
@@ -190,10 +192,10 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
 
 	int i, k;
         uint32_t j;
-	const int batchsize = ITER >> ( 2 + bfactor );
+	const int batchsize = (ITERATIONS * 2) >> ( 2 + bfactor );
 	const int start = partidx * batchsize;
 	const int end = start + batchsize;
-	uint32_t * long_state = &d_long_state[(IndexType) thread << 19];
+	uint32_t * long_state = &d_long_state[(IndexType) thread << THREAD_SHIFT];
 	uint32_t * ctx_a = d_ctx_a + thread * 4;
 	uint32_t * ctx_b = d_ctx_b + thread * 4;
 	uint32_t a, d[2];
@@ -207,7 +209,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
 		#pragma unroll 2
 		for ( int x = 0; x < 2; ++x )
 		{
-			j = ( ( shuffle(sPtr,sub, a, 0) & 0x1FFFF0 ) >> 2 ) + sub;
+			j = ( ( shuffle(sPtr,sub, a, 0) & MASK ) >> 2 ) + sub;
 
 			const uint32_t x_0 = loadGlobal32<uint32_t>( long_state + j );
 			const uint32_t x_1 = shuffle(sPtr,sub, x_0, sub + 1);
@@ -225,8 +227,8 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
 			//long_state[j] = d[0] ^ d[1];
 			storeGlobal32( long_state + j, d[0] ^ d[1] );
 
-			//MUL_SUM_XOR_DST(c, a, &long_state[((uint32_t *)c)[0] & 0x1FFFF0]);
-			j = ( ( *t1 & 0x1FFFF0 ) >> 2 ) + sub;
+			//MUL_SUM_XOR_DST(c, a, &long_state[((uint32_t *)c)[0] & MASK]);
+			j = ( ( *t1 & MASK ) >> 2 ) + sub;
 
 			uint32_t yy[2];
 			*( (uint64_t*) yy ) = loadGlobal64<uint64_t>( ( (uint64_t *) long_state )+( j >> 1 ) );
@@ -255,6 +257,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
 	}
 }
 
+template<size_t ITERATIONS, size_t THREAD_SHIFT>
 __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int partidx, const uint32_t * __restrict__ long_state, uint32_t * __restrict__ d_ctx_state, uint32_t * __restrict__ d_ctx_key2 )
 {
 	__shared__ uint32_t sharedMemory[1024];
@@ -265,7 +268,7 @@ __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int parti
 	int thread = ( blockDim.x * blockIdx.x + threadIdx.x ) >> 3;
 	int sub = ( threadIdx.x & 7 ) << 2;
 
-	const int batchsize = 0x80000 >> bfactor;
+	const int batchsize = ITERATIONS >> bfactor;
 	const int start = partidx * batchsize;
 	const int end = start + batchsize;
 
@@ -281,7 +284,7 @@ __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int parti
 	{
 #pragma unroll
 		for ( int j = 0; j < 4; ++j )
-			text[j] ^= long_state[((IndexType) thread << 19) + (sub + i + j)];
+			text[j] ^= long_state[((IndexType) thread << THREAD_SHIFT) + (sub + i + j)];
 
 		cn_aes_pseudo_round_mut( sharedMemory, text, key );
 	}
@@ -289,7 +292,8 @@ __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int parti
 	MEMCPY8( d_ctx_state + thread * 50 + sub + 16, text, 2 );
 }
 
-extern "C" void cryptonight_core_cpu_hash(nvid_ctx* ctx)
+template<size_t ITERATIONS, size_t THREAD_SHIFT, uint32_t MASK>
+void cryptonight_core_cpu_hash(nvid_ctx* ctx)
 {
 	dim3 grid( ctx->device_blocks );
 	dim3 block( ctx->device_threads );
@@ -311,7 +315,7 @@ extern "C" void cryptonight_core_cpu_hash(nvid_ctx* ctx)
 
 	for ( int i = 0; i < partcountOneThree; i++ )
 	{
-		CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase1<<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads,
+		CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase1<ITERATIONS,THREAD_SHIFT><<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads,
 			bfactorOneThree, i,
 			ctx->d_long_state, ctx->d_ctx_state, ctx->d_ctx_key1 ));
 
@@ -321,7 +325,7 @@ extern "C" void cryptonight_core_cpu_hash(nvid_ctx* ctx)
 
 	for ( int i = 0; i < partcount; i++ )
 	{
-        CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase2<<<
+        CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase2<ITERATIONS,THREAD_SHIFT,MASK><<<
             grid,
             block4,
             block4.x * sizeof(uint32_t) * static_cast< int >( ctx->device_arch[0] < 3 )
@@ -339,7 +343,7 @@ extern "C" void cryptonight_core_cpu_hash(nvid_ctx* ctx)
 
 	for ( int i = 0; i < partcountOneThree; i++ )
 	{
-		CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase3<<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads,
+		CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase3<ITERATIONS,THREAD_SHIFT><<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads,
 			bfactorOneThree, i,
 			ctx->d_long_state,
 			ctx->d_ctx_state, ctx->d_ctx_key2 ));
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
index 7734473..abca489 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
@@ -5,6 +5,7 @@
 #include <cuda_runtime.h>
 #include <device_functions.hpp>
 #include  <algorithm>
+#include "xmrstak/jconf.hpp"
 
 #ifdef __CUDACC__
 __constant__
@@ -188,8 +189,18 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx)
 	cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
 	cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);
 
+	size_t hashMemSize;
+	if(::jconf::inst()->IsCurrencyMonero())
+	{
+		hashMemSize = MONERO_MEMORY;
+	}
+	else
+	{
+		hashMemSize = AEON_MEMORY;
+	}
+
 	size_t wsize = ctx->device_blocks * ctx->device_threads;
-	CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_long_state, (size_t)MEMORY * wsize));
+	CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_long_state, hashMemSize * wsize));
 	CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_state, 50 * sizeof(uint32_t) * wsize));
 	CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_key1, 40 * sizeof(uint32_t) * wsize));
 	CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_key2, 40 * sizeof(uint32_t) * wsize));
@@ -343,13 +354,23 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
 		ctx->total_device_memory = totalMemory;
 		ctx->free_device_memory = freeMemory;
 
+		size_t hashMemSize;
+		if(::jconf::inst()->IsCurrencyMonero())
+		{
+			hashMemSize = MONERO_MEMORY;
+		}
+		else
+		{
+			hashMemSize = AEON_MEMORY;
+		}
+
 		// keep 64MiB memory free (value is randomly chosen)
 		// 200byte are meta data memory (result nonce, ...)
 		size_t availableMem = freeMemory - (64u * 1024 * 1024) - 200u;
 		size_t limitedMemory = std::min(availableMem, maxMemUsage);
 		// up to 920bytes extra memory is used per thread for some kernel (lmem/local memory)
 		// 680bytes are extra meta data memory per hash
-		size_t perThread = size_t(MEMORY) + 740u + 680u;
+		size_t perThread = hashMemSize + 740u + 680u;
 		size_t max_intensity = limitedMemory / perThread;
 		ctx->device_threads = max_intensity / ctx->device_blocks;
 		// use only odd number of threads
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.hpp b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.hpp
index 3ccdcd6..055a8bd 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.hpp
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.hpp
@@ -1,5 +1,7 @@
 #pragma once
 
+#include "xmrstak/backend/cryptonight.hpp"
+
 #ifdef __INTELLISENSE__
 #define __CUDA_ARCH__ 520
 /* avoid red underlining */
@@ -18,8 +20,6 @@ struct uint3  blockDim;
 #define __shfl(a,b,c) 1
 #endif
 
-#define MEMORY         (1 << 21) // 2 MiB / 2097152 B
-#define ITER           (1 << 20) // 1048576
 #define AES_BLOCK_SIZE  16
 #define AES_KEY_SIZE    32
 #define INIT_SIZE_BLK   8