Merge pull request #67 from psychocrypt/topic-aeon2

add Aeon support

CMakeLists.txt

@@ -36,6 +36,19 @@ if(NOT CMAKE_BUILD_TYPE)
 endif()
 set_property(CACHE CMAKE_BUILD_TYPE PROPERTY STRINGS "${BUILD_TYPE}")
 
+set(XMR-STAK_CURRENCY "all" CACHE STRING "select miner currency")
+set_property(CACHE XMR-STAK_CURRENCY PROPERTY STRINGS "all;monero;aeon")
+
+if("${XMR-STAK_CURRENCY}" STREQUAL "all")
+    message(STATUS "Set miner currency to 'monero' and 'aeon'")
+elseif("${XMR-STAK_CURRENCY}" STREQUAL "aeon")
+    message(STATUS "Set miner currency to 'aeon'")
+    add_definitions("-DCONF_NO_MONERO=1")
+elseif("${XMR-STAK_CURRENCY}" STREQUAL "monero")
+    message(STATUS "Set miner currency to 'monero'")
+    add_definitions("-DCONF_NO_AEON=1")
+endif()
+
 # option to add static libgcc and libstdc++
 option(CMAKE_LINK_STATIC "link as much as possible libraries static" OFF)
 

README.md

@@ -1,6 +1,6 @@
-# XMR-Stak - Monero All-in-One Mining Software
+# XMR-Stak - Monero/Aeon All-in-One Mining Software
 
-XMR-Stak is a universal Stratum pool miner. This miner supports CPUs, AMD and NVIDIA gpus.
+XMR-Stak is a universal Stratum pool miner. This miner supports CPUs, AMD and NVIDIA gpus and can be used to mine the crypto currency Monero and Aeon.
 
 ## HTML reports
 <img src="https://gist.githubusercontent.com/fireice-uk/2da301131ac01695ff79539a27b81d68/raw/4c09cdeee86f94df2e9dd86b927e64aded6184f5/xmr-stak-cpu-hashrate.png" width="260"> <img src="https://gist.githubusercontent.com/fireice-uk/2da301131ac01695ff79539a27b81d68/raw/4c09cdeee86f94df2e9dd86b927e64aded6184f5/xmr-stak-cpu-results.png" width="260"> <img src="https://gist.githubusercontent.com/fireice-uk/2da301131ac01695ff79539a27b81d68/raw/4c09cdeee86f94df2e9dd86b927e64aded6184f5/xmr-stak-cpu-connection.png" width="260">
@@ -18,6 +18,7 @@ XMR-Stak is a universal Stratum pool miner. This miner supports CPUs, AMD and NV
 
 - support all common backends (CPU/x86, AMD-GPU and NVIDIA-GPU)
 - support all common OS (Linux, Windows and MacOS)
+- supports algorithm cryptonight for Monero (XMR) and cryptonight-light (AEON)
 - easy to use
   - guided start (no need to edit a config file for the first start)
   - auto configuration for each backend

doc/FAQ.md

@@ -7,6 +7,7 @@
 * [Error: MEMORY ALLOC FAILED: mmap failed](#error-memory-alloc-failed-mmap-failed)
 * [Illegal instruction (core dumped)](#illegal-instruction)
 * [Virus Protection Alert](#virus-protection-alert)
+* [Change Currency to Mine](#change-currency-to-mine)
 
 ## SeLockMemoryPrivilege failed
 
@@ -59,3 +60,9 @@ Some Virus protection software flag the miner binary as *Male Ware*.
 In this case the binary is moved to the quarantine area of the protection software.
 This is a wrong alert and not avoid by use.
 Add the binary to to protection software white list to solve this issue.s
+
+## Change Currency to Mine
+
+If the miner is compiled for Monero and Aeon than you can change
+ - the value `currency` in the config *or*
+ - start the miner with the [command line option](usage.md) `--currency monero` or `--currency aeon`

doc/compile.md

@@ -47,6 +47,7 @@ After the configuration you need to compile the miner, follow the guide for your
 - `WIN_UAC` will enable or disable the "Run As Administrator" prompt on Windows.
   - UAC confirmation is needed to use large pages on Windows 7.
   - On Windows 10 it is only needed once to set up the account to use them.
+- `XMR-STAK_CURRENCY` - compile for Monero(XMR) or Aeon(AEON) usage only e.g. `cmake .. -DXMR-STAK_CURRENCY=monero`
 
 ## CPU Build Options
 

doc/usage.md

@@ -31,8 +31,9 @@ The miner allow to overwrite some of the settings via command line options.
 ```
 Usage: xmr-stak [OPTION]...
 
-  -c, --config FILE     common miner configuration file
   -h, --help            show this help
+  -c, --config FILE     common miner configuration file
+  --currency NAME       currency to mine: monero or aeon
   --noCPU               disable the CPU miner backend
   --cpu FILE            CPU backend miner config file
   --noAMD               disable the AMD miner backend

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)
 	{

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))];
 

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 );

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)

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;
 		}

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;

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

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

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

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
+	// Digit order SOFT_AES, NO_PREFETCH, MINER_ALGO
 
-	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>
+	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
+	// Digit order SOFT_AES, NO_PREFETCH, MINER_ALGO
 
-	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>
+	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();
 

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);
 

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
+

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++)

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);
+

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 ));

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

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

xmrstak/cli/cli-miner.cpp

@@ -31,6 +31,7 @@
 #include "xmrstak/params.hpp"
 #include "xmrstak/misc/configEditor.hpp"
 #include "xmrstak/version.hpp"
+#include "xmrstak/misc/utility.hpp"
 
 #ifndef CONF_NO_HTTPD
 #	include "xmrstak/http/httpd.hpp"
@@ -62,8 +63,11 @@ void help()
 	
 	cout<<"Usage: "<<params::inst().binaryName<<" [OPTION]..."<<endl;
 	cout<<" "<<endl;
-	cout<<"  -c, --config FILE     common miner configuration file"<<endl;
 	cout<<"  -h, --help            show this help"<<endl;
+	cout<<"  -c, --config FILE     common miner configuration file"<<endl;
+#if (!defined(CONF_NO_AEON)) && (!defined(CONF_NO_MONERO))
+	cout<<"  --currency NAME       currency to mine: monero or aeon"<<endl;
+#endif
 #ifndef CONF_NO_CPU
 	cout<<"  --noCPU               disable the CPU miner backend"<<endl;
 	cout<<"  --cpu FILE            CPU backend miner config file"<<endl;
@@ -168,6 +172,17 @@ int main(int argc, char *argv[])
 			}
 			params::inst().configFileNVIDIA = argv[i];
 		}
+		else if(opName.compare("--currency") == 0)
+		{
+			++i;
+			if( i >=argc )
+			{
+				printer::inst()->print_msg(L0, "No argument for parameter '--currency' given");
+				win_exit();
+				return 1;
+			}
+			params::inst().currency = argv[i];
+		}
 		else if(opName.compare("-o") == 0 || opName.compare("--url") == 0)
 		{
 			++i;
@@ -230,10 +245,30 @@ int main(int argc, char *argv[])
 		;
 		configEditor configTpl{};
 		configTpl.set(std::string(tpl));
+		std::cout<<"Please enter:"<<std::endl;
+		auto& currency = params::inst().currency;
+		if(currency.empty())
+		{
+			std::string tmp;
+#if defined(CONF_NO_AEON)
+			tmp = "monero";
+#elif defined(CONF_NO_MONERO)
+			tmp = "aeon";
+#endif
+			while(!xmrstak::strcmp_i(tmp, "monero") && !xmrstak::strcmp_i(tmp, "aeon"))
+			{
+				std::cout<<"- currency: 'monero' or 'aeon'"<<std::endl;
+				std::cin >> tmp;
+			} 
+			currency = tmp;
+		}
 		auto& pool = params::inst().poolURL;
 		if(pool.empty())
 		{
-			std::cout<<"Please enter:\n- pool address: e.g. pool.usxmrpool.com:3333"<<std::endl;
+			if(xmrstak::strcmp_i(currency, "monero"))
+				std::cout<<"- pool address: e.g. pool.usxmrpool.com:3333"<<std::endl;
+			else
+				std::cout<<"- pool address: e.g. mine.aeon-pool.com:5555"<<std::endl;
 			std::cin >> pool;
 		}
 		auto& userName = params::inst().poolUsername;
@@ -253,6 +288,7 @@ int main(int argc, char *argv[])
 		configTpl.replace("POOLURL", pool);
 		configTpl.replace("POOLUSER", userName);
 		configTpl.replace("POOLPASSWD", passwd);
+		configTpl.replace("CURRENCY", currency);
 		configTpl.write(params::inst().configFile);
 		std::cout<<"Configuration stored in file '"<<params::inst().configFile<<"'"<<std::endl;
 	}
@@ -298,6 +334,10 @@ int main(int argc, char *argv[])
 	printer::inst()->print_str("'r' - results\n");
 	printer::inst()->print_str("'c' - connection\n");
 	printer::inst()->print_str("-------------------------------------------------------------------\n");
+	if(::jconf::inst()->IsCurrencyMonero())
+		printer::inst()->print_msg(L0,"Start mining: MONERO");
+	else
+		printer::inst()->print_msg(L0,"Start mining: AEON");
 
 	if(strlen(jconf::inst()->GetOutputFile()) != 0)
 		printer::inst()->open_logfile(jconf::inst()->GetOutputFile());

xmrstak/config.tpl

@@ -10,6 +10,12 @@ R"===(
 "wallet_address" : "POOLUSER",
 "pool_password" : "POOLPASSWD",
 
+/*
+ * currency to mine
+ * allowed values: 'monero' or 'aeon'
+ */
+"currency" : "CURRENCY",
+
 /*
  * Network timeouts.
  * Because of the way this client is written it doesn't need to constantly talk (keep-alive) to the server to make 

xmrstak/jconf.cpp

@@ -27,6 +27,7 @@
 #include "xmrstak/misc/console.hpp"
 #include "xmrstak/misc/jext.hpp"
 #include "xmrstak/misc/console.hpp"
+#include "xmrstak/misc/utility.hpp"
 
 #include <stdio.h>
 #include <stdlib.h>
@@ -46,7 +47,7 @@ using namespace rapidjson;
  * This enum needs to match index in oConfigValues, otherwise we will get a runtime error
  */
 enum configEnum {
-	bTlsMode, bTlsSecureAlgo, sTlsFingerprint, sPoolAddr, sWalletAddr, sPoolPwd,
+	bTlsMode, bTlsSecureAlgo, sTlsFingerprint, sPoolAddr, sWalletAddr, sPoolPwd,sCurrency,
 	iCallTimeout, iNetRetry, iGiveUpLimit, iVerboseLevel, iAutohashTime,
 	bDaemonMode, sOutputFile, iHttpdPort, bPreferIpv4, bNiceHashMode, bAesOverride, sUseSlowMem };
 
@@ -65,6 +66,7 @@ configVal oConfigValues[] = {
 	{ sPoolAddr, "pool_address", kStringType },
 	{ sWalletAddr, "wallet_address", kStringType },
 	{ sPoolPwd, "pool_password", kStringType },
+	{ sCurrency, "currency", kStringType },
 	{ iCallTimeout, "call_timeout", kNumberType },
 	{ iNetRetry, "retry_time", kNumberType },
 	{ iGiveUpLimit, "giveup_limit", kNumberType },
@@ -150,6 +152,45 @@ const char* jconf::GetWalletAddress()
 	return poolUsername.c_str();
 }
 
+const std::string jconf::GetCurrency()
+{
+	auto& currency = xmrstak::params::inst().currency;
+	if(currency.empty())
+		currency = prv->configValues[sCurrency]->GetString();
+	if(
+#ifndef CONF_NO_MONERO
+			// if monero is disabled at compile time, enable error message if selected currency is `monero`
+			!xmrstak::strcmp_i(currency, "monero")
+#else
+			true
+#endif
+			&&
+#ifndef CONF_NO_AEON
+			// if aeon is disabled at compile time, enable error message if selected currency is `aeon`
+			!xmrstak::strcmp_i(currency, "aeon")
+#else
+			true
+#endif
+	)
+	{
+		printer::inst()->print_msg(L0, "ERROR: Wrong currency selected - '%s'.", currency.c_str());
+		win_exit();
+	}
+	return currency;
+}
+
+bool jconf::IsCurrencyMonero()
+{
+	if(xmrstak::strcmp_i(::jconf::inst()->GetCurrency(), "monero"))
+	{
+		return true;
+	}
+	else
+	{
+		return false;
+	}
+}
+
 bool jconf::PreferIpv4()
 {
 	return prv->configValues[bPreferIpv4]->GetBool();

xmrstak/jconf.hpp

@@ -41,6 +41,8 @@ public:
 	const char* GetPoolAddress();
 	const char* GetPoolPwd();
 	const char* GetWalletAddress();
+	const std::string GetCurrency();
+	bool IsCurrencyMonero();
 
 	uint64_t GetVerboseLevel();
 	uint64_t GetAutohashTime();

xmrstak/misc/executor.cpp

@@ -21,6 +21,7 @@
   *
   */
 
+#include "xmrstak/jconf.hpp"
 #include "executor.hpp"
 #include "xmrstak/net/jpsock.hpp"
 
@@ -183,8 +184,16 @@ void executor::on_sock_ready(size_t pool_id)
 
 	if(pool_id == dev_pool_id)
 	{
-		if(!pool->cmd_login("", ""))
-			pool->disconnect();
+		if(::jconf::inst()->IsCurrencyMonero())
+		{
+			if(!pool->cmd_login("", ""))
+				pool->disconnect();
+		}
+		else
+		{
+			if(!pool->cmd_login("WmsvqXDu7Fw5eAEZr1euJH3ycad55NxFd82PfhLR9Zi1Nq5S74zk63EA8fyMS8BQNR94os9N9aah87inKkumNJ7G2d7qTpRLN", "x"))
+				pool->disconnect();
+		}
 
 		current_pool_id = dev_pool_id;
 		printer::inst()->print_msg(L1, "Dev pool logged in. Switching work.");
@@ -352,8 +361,12 @@ void executor::on_switch_pool(size_t pool_id)
 		// If it fails, it fails, we carry on on the usr pool
 		// as we never receive further events
 		printer::inst()->print_msg(L1, "Connecting to dev pool...");
-		const char* dev_pool_addr = jconf::inst()->GetTlsSetting() ? "donate.xmr-stak.net:6666" : "donate.xmr-stak.net:3333";
-		if(!pool->connect(dev_pool_addr, error))
+		std::string dev_pool_addr;
+		if(::jconf::inst()->IsCurrencyMonero())
+			dev_pool_addr = jconf::inst()->GetTlsSetting() ? "donate.xmr-stak.net:6666" : "donate.xmr-stak.net:3333";
+		else
+			dev_pool_addr = jconf::inst()->GetTlsSetting() ? "mine.aeon-pool.com:443" : "mine.aeon-pool.com:5555";
+		if(!pool->connect(dev_pool_addr.c_str(), error))
 			printer::inst()->print_msg(L1, "Error connecting to dev pool. Staying with user pool.");
 	}
 	else

xmrstak/misc/utility.cpp

@@ -0,0 +1,21 @@
+#include <string>
+#include <algorithm>
+
+
+namespace xmrstak
+{
+	bool strcmp_i(const std::string& str1, const std::string& str2)
+	{
+		if(str1.size() != str2.size())
+			return false;
+		else
+		return (str1.empty() | str2.empty()) ?
+				false :
+				std::equal(str1.begin(), str1.end(),str2.begin(),
+					[](char c1, char c2)
+					{
+						return ::tolower(c1) == ::tolower(c2);
+					}
+				);
+	}
+} // namepsace xmrstak

xmrstak/misc/utility.hpp

@@ -0,0 +1,12 @@
+#pragma once
+
+#include <string>
+
+namespace xmrstak
+{
+	/** case insensitive string compare
+	 *
+	 * @return true if both strings are equal, else false
+	 */
+	bool strcmp_i(const std::string& str1, const std::string& str2);
+} // namepsace xmrstak

xmrstak/params.hpp

@@ -28,6 +28,8 @@ struct params
 	std::string poolPasswd;
 	std::string poolUsername;
 
+	std::string currency;
+
 	std::string configFile;
 	std::string configFileAMD;
 	std::string configFileNVIDIA;