summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
authorfireice-uk <fireice-uk@users.noreply.github.com>2017-10-27 20:37:12 +0100
committerGitHub <noreply@github.com>2017-10-27 20:37:12 +0100
commitc192b97b67bf648540bf058f21b7873bfe886a82 (patch)
treefe71edd305f523021c9038dab5595551c7e4c7e2
parent91b307859ea97cc1abdb17da01b94919d3521803 (diff)
parent0117ed609e69af977fabaf9c28bbf8cf1300ee89 (diff)
downloadxmr-stak-c192b97b67bf648540bf058f21b7873bfe886a82.zip
xmr-stak-c192b97b67bf648540bf058f21b7873bfe886a82.tar.gz
Merge pull request #67 from psychocrypt/topic-aeon2
add Aeon support
-rw-r--r--CMakeLists.txt13
-rw-r--r--README.md5
-rw-r--r--doc/FAQ.md7
-rw-r--r--doc/compile.md1
-rw-r--r--doc/usage.md3
-rw-r--r--xmrstak/backend/amd/amd_gpu/gpu.cpp26
-rw-r--r--xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl22
-rw-r--r--xmrstak/backend/amd/autoAdjust.hpp14
-rw-r--r--xmrstak/backend/amd/minethd.cpp2
-rw-r--r--xmrstak/backend/cpu/autoAdjust.hpp26
-rw-r--r--xmrstak/backend/cpu/autoAdjustHwloc.hpp21
-rw-r--r--xmrstak/backend/cpu/crypto/cryptonight.h2
-rw-r--r--xmrstak/backend/cpu/crypto/cryptonight_aesni.h52
-rw-r--r--xmrstak/backend/cpu/crypto/cryptonight_common.cpp38
-rw-r--r--xmrstak/backend/cpu/minethd.cpp126
-rw-r--r--xmrstak/backend/cpu/minethd.hpp4
-rw-r--r--xmrstak/backend/cryptonight.hpp12
-rw-r--r--xmrstak/backend/nvidia/minethd.cpp20
-rw-r--r--xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp5
-rw-r--r--xmrstak/backend/nvidia/nvcc_code/cuda_core.cu32
-rw-r--r--xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu25
-rw-r--r--xmrstak/backend/nvidia/nvcc_code/cuda_extra.hpp4
-rw-r--r--xmrstak/cli/cli-miner.cpp44
-rw-r--r--xmrstak/config.tpl6
-rw-r--r--xmrstak/jconf.cpp43
-rw-r--r--xmrstak/jconf.hpp2
-rw-r--r--xmrstak/misc/executor.cpp21
-rw-r--r--xmrstak/misc/utility.cpp21
-rw-r--r--xmrstak/misc/utility.hpp12
-rw-r--r--xmrstak/params.hpp2
30 files changed, 477 insertions, 134 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 1cc2ef6..595631d 100644
--- a/CMakeLists.txt
+++ b/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)
diff --git a/README.md b/README.md
index cea9cd4..15f641c 100644
--- a/README.md
+++ b/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
diff --git a/doc/FAQ.md b/doc/FAQ.md
index 995a421..215048d 100644
--- a/doc/FAQ.md
+++ b/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`
diff --git a/doc/compile.md b/doc/compile.md
index 248d6d9..537a736 100644
--- a/doc/compile.md
+++ b/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
diff --git a/doc/usage.md b/doc/usage.md
index 42a29fd..9a38384 100644
--- a/doc/usage.md
+++ b/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
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
diff --git a/xmrstak/cli/cli-miner.cpp b/xmrstak/cli/cli-miner.cpp
index 8939803..f571ad6 100644
--- a/xmrstak/cli/cli-miner.cpp
+++ b/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());
diff --git a/xmrstak/config.tpl b/xmrstak/config.tpl
index 60d85cd..032d483 100644
--- a/xmrstak/config.tpl
+++ b/xmrstak/config.tpl
@@ -11,6 +11,12 @@ R"===(
"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
* sure it is there. We detect a buggy / overloaded server by the call timeout. The default values will be ok for
diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp
index c033e66..4b23ed5 100644
--- a/xmrstak/jconf.cpp
+++ b/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();
diff --git a/xmrstak/jconf.hpp b/xmrstak/jconf.hpp
index 688ffe1..48b47b5 100644
--- a/xmrstak/jconf.hpp
+++ b/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();
diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp
index 7fc46e4..b469dc2 100644
--- a/xmrstak/misc/executor.cpp
+++ b/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
diff --git a/xmrstak/misc/utility.cpp b/xmrstak/misc/utility.cpp
new file mode 100644
index 0000000..3b1369a
--- /dev/null
+++ b/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
diff --git a/xmrstak/misc/utility.hpp b/xmrstak/misc/utility.hpp
new file mode 100644
index 0000000..b2e841d
--- /dev/null
+++ b/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
diff --git a/xmrstak/params.hpp b/xmrstak/params.hpp
index dddf82e..6127212 100644
--- a/xmrstak/params.hpp
+++ b/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;
OpenPOWER on IntegriCloud