From 5a2c4444a4e2b5ddc5e5dbf98990dcad45e32fe8 Mon Sep 17 00:00:00 2001
From: psychocrypt <psychocryptHPC@gmail.com>
Date: Tue, 27 Mar 2018 20:54:47 +0200
Subject: [PATCH] POW AEON v7
- add new pow for AEON
- fix missing cryptonight-heavy selection for multi hashes
---
xmrstak/backend/amd/amd_gpu/gpu.cpp | 6 +-
xmrstak/backend/amd/minethd.cpp | 33 +++---
.../backend/cpu/crypto/cryptonight_aesni.h | 30 +++---
xmrstak/backend/cpu/minethd.cpp | 100 ++++++++++++------
xmrstak/backend/cryptonight.hpp | 18 +++-
xmrstak/backend/nvidia/minethd.cpp | 34 +++---
xmrstak/backend/nvidia/nvcc_code/cuda_core.cu | 13 ++-
xmrstak/jconf.cpp | 22 ++--
xmrstak/jconf.hpp | 3 +
xmrstak/misc/executor.cpp | 7 ++
xmrstak/net/jpsock.cpp | 3 +
11 files changed, 167 insertions(+), 102 deletions(-)
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp
index 8d0fd32..80d852a 100644
--- a/xmrstak/backend/amd/amd_gpu/gpu.cpp
+++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp
@@ -936,7 +936,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
/// @todo only activate if currency is monero
int cn_kernel_offset = 0;
- if(miner_algo == cryptonight_monero && version >= 7)
+ if((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon ) && version >= 7)
{
cn_kernel_offset = 6;
}
@@ -962,7 +962,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
return(ERR_OCL_API);
}
- if(miner_algo == cryptonight_monero && version >= 7)
+ if((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon ) && version >= 7)
{
// Input
if ((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS)
@@ -1130,7 +1130,7 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo,
size_t tmpNonce = ctx->Nonce;
/// @todo only activate if currency is monero
int cn_kernel_offset = 0;
- if(miner_algo == cryptonight_monero && version >= 7)
+ if((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version >= 7)
{
cn_kernel_offset = 6;
}
diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp
index 46a04d5..78774c0 100644
--- a/xmrstak/backend/amd/minethd.cpp
+++ b/xmrstak/backend/amd/minethd.cpp
@@ -192,14 +192,9 @@ void minethd::work_main()
cryptonight_ctx* cpu_ctx;
cpu_ctx = cpu::minethd::minethd_alloc_ctx();
auto miner_algo = ::jconf::inst()->GetMiningAlgo();
- cn_hash_fun hash_fun;
- if(miner_algo == cryptonight_monero || miner_algo == cryptonight_heavy)
- {
- // start with cryptonight and switch later if fork version is reached
- hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight);
- }
- else
- hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
+
+ // start with root algorithm and switch later if fork version is reached
+ cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, ::jconf::inst()->GetMiningAlgoRoot());
globalStates::inst().iConsumeCnt++;
@@ -219,13 +214,12 @@ void minethd::work_main()
consume_work();
uint8_t new_version = oWork.getVersion();
- if(miner_algo == cryptonight_monero && version < 7 && new_version >= 7)
- {
- hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_monero);
- }
- else if(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3)
+ const bool time_to_fork =
+ ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version < 7 && new_version >= 7) ||
+ (miner_algo == cryptonight_heavy && version < 3 && new_version >= 3);
+ if(time_to_fork)
{
- hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy);
+ hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
}
version = new_version;
continue;
@@ -281,13 +275,12 @@ void minethd::work_main()
consume_work();
uint8_t new_version = oWork.getVersion();
- if(miner_algo == cryptonight_monero && version < 7 && new_version >= 7)
- {
- hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_monero);
- }
- else if(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3)
+ const bool time_to_fork =
+ ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version < 7 && new_version >= 7) ||
+ (miner_algo == cryptonight_heavy && version < 3 && new_version >= 3);
+ if(time_to_fork)
{
- hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy);
+ hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
}
version = new_version;
}
diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
index 85373e8..5203de8 100644
--- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
+++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
@@ -444,7 +444,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
constexpr size_t MEM = cn_select_memory<ALGO>();
- if(ALGO == cryptonight_monero && len < 43)
+ if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43)
{
memset(output, 0, 32);
return;
@@ -453,7 +453,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
keccak((const uint8_t *)input, len, ctx0->hash_state, 200);
uint64_t monero_const;
- if(ALGO == cryptonight_monero)
+ if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
{
monero_const = *reinterpret_cast<const uint64_t*>(reinterpret_cast<const uint8_t*>(input) + 35);
monero_const ^= *(reinterpret_cast<const uint64_t*>(ctx0->hash_state) + 24);
@@ -482,7 +482,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
else
cx = _mm_aesenc_si128(cx, _mm_set_epi64x(ah0, al0));
- if(ALGO == cryptonight_monero)
+ if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
cryptonight_monero_tweak((uint64_t*)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
else
_mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
@@ -506,7 +506,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
_mm_prefetch((const char*)&l0[al0 & MASK], _MM_HINT_T0);
ah0 += lo;
- if(ALGO == cryptonight_monero)
+ if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
((uint64_t*)&l0[idx0 & MASK])[1] = ah0 ^ monero_const;
else
((uint64_t*)&l0[idx0 & MASK])[1] = ah0;
@@ -544,7 +544,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
constexpr size_t MEM = cn_select_memory<ALGO>();
- if(ALGO == cryptonight_monero && len < 43)
+ if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43)
{
memset(output, 0, 64);
return;
@@ -554,7 +554,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
keccak((const uint8_t *)input+len, len, ctx[1]->hash_state, 200);
uint64_t monero_const_0, monero_const_1;
- if(ALGO == cryptonight_monero)
+ if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
{
monero_const_0 = *reinterpret_cast<const uint64_t*>(reinterpret_cast<const uint8_t*>(input) + 35);
monero_const_0 ^= *(reinterpret_cast<const uint64_t*>(ctx[0]->hash_state) + 24);
@@ -592,7 +592,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
else
cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh0, axl0));
- if(ALGO == cryptonight_monero)
+ if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
cryptonight_monero_tweak((uint64_t*)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
else
_mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
@@ -610,7 +610,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
else
cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh1, axl1));
- if(ALGO == cryptonight_monero)
+ if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
cryptonight_monero_tweak((uint64_t*)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx));
else
_mm_store_si128((__m128i *)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx));
@@ -631,7 +631,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
axh0 += lo;
((uint64_t*)&l0[idx0 & MASK])[0] = axl0;
- if(ALGO == cryptonight_monero)
+ if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
((uint64_t*)&l0[idx0 & MASK])[1] = axh0 ^ monero_const_0;
else
((uint64_t*)&l0[idx0 & MASK])[1] = axh0;
@@ -662,7 +662,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
axh1 += lo;
((uint64_t*)&l1[idx1 & MASK])[0] = axl1;
- if(ALGO == cryptonight_monero)
+ if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
((uint64_t*)&l1[idx1 & MASK])[1] = axh1 ^ monero_const_1;
else
((uint64_t*)&l1[idx1 & MASK])[1] = axh1;
@@ -709,7 +709,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
else \
c = _mm_aesenc_si128(c, a); \
b = _mm_xor_si128(b, c); \
- if(ALGO == cryptonight_monero) \
+ if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) \
cryptonight_monero_tweak((uint64_t*)ptr, b); \
else \
_mm_store_si128(ptr, b);\
@@ -724,7 +724,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
#define CN_STEP4(a, b, c, l, mc, ptr, idx) \
lo = _umul128(idx, _mm_cvtsi128_si64(b), &hi); \
a = _mm_add_epi64(a, _mm_set_epi64x(lo, hi)); \
- if(ALGO == cryptonight_monero) \
+ if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) \
_mm_store_si128(ptr, _mm_xor_si128(a, mc)); \
else \
_mm_store_si128(ptr, a);\
@@ -751,7 +751,7 @@ void cryptonight_triple_hash(const void* input, size_t len, void* output, crypto
constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
constexpr size_t MEM = cn_select_memory<ALGO>();
- if(ALGO == cryptonight_monero && len < 43)
+ if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43)
{
memset(output, 0, 32 * 3);
return;
@@ -845,7 +845,7 @@ void cryptonight_quad_hash(const void* input, size_t len, void* output, cryptoni
constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
constexpr size_t MEM = cn_select_memory<ALGO>();
- if(ALGO == cryptonight_monero && len < 43)
+ if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43)
{
memset(output, 0, 32 * 4);
return;
@@ -954,7 +954,7 @@ void cryptonight_penta_hash(const void* input, size_t len, void* output, crypton
constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
constexpr size_t MEM = cn_select_memory<ALGO>();
- if(ALGO == cryptonight_monero && len < 43)
+ if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43)
{
memset(output, 0, 32 * 5);
return;
diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp
index e263aca..d7e242b 100644
--- a/xmrstak/backend/cpu/minethd.cpp
+++ b/xmrstak/backend/cpu/minethd.cpp
@@ -282,6 +282,9 @@ bool minethd::self_test()
else if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero)
{
}
+ else if(::jconf::inst()->GetMiningAlgo() == cryptonight_aeon)
+ {
+ }
for (int i = 0; i < MAX_N; i++)
cryptonight_free_ctx(ctx[i]);
@@ -366,6 +369,9 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr
case cryptonight_heavy:
algv = 3;
break;
+ case cryptonight_aeon:
+ algv = 4;
+ break;
default:
algv = 2;
break;
@@ -387,7 +393,11 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr
cryptonight_hash<cryptonight_heavy, false, false>,
cryptonight_hash<cryptonight_heavy, true, false>,
cryptonight_hash<cryptonight_heavy, false, true>,
- cryptonight_hash<cryptonight_heavy, true, true>
+ cryptonight_hash<cryptonight_heavy, true, true>,
+ cryptonight_hash<cryptonight_aeon, false, false>,
+ cryptonight_hash<cryptonight_aeon, true, false>,
+ cryptonight_hash<cryptonight_aeon, false, true>,
+ cryptonight_hash<cryptonight_aeon, true, true>
};
std::bitset<2> digit;
@@ -447,21 +457,15 @@ void minethd::work_main()
if(oWork.bNiceHash)
result.iNonce = *piNonce;
- if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero)
- {
- if(oWork.bWorkBlob[0] >= 7)
- hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight_monero);
- else
- hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight);
- }
+ auto miner_algo = ::jconf::inst()->GetMiningAlgo();
+ const bool time_to_fork =
+ ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && oWork.bWorkBlob[0] >= 7) ||
+ (miner_algo == cryptonight_heavy && oWork.bWorkBlob[0] >= 3);
- if(::jconf::inst()->GetMiningAlgo() == cryptonight_heavy)
- {
- if(oWork.bWorkBlob[0] >= 3)
- hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight_heavy);
- else
- hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight);
- }
+ if(time_to_fork)
+ hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
+ else
+ hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgoRoot());
while(globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo)
{
@@ -512,6 +516,12 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
case cryptonight_monero:
algv = 0;
break;
+ case cryptonight_heavy:
+ algv = 3;
+ break;
+ case cryptonight_aeon:
+ algv = 4;
+ break;
default:
algv = 2;
break;
@@ -534,6 +544,7 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
cryptonight_penta_hash<cryptonight_monero, true, false>,
cryptonight_penta_hash<cryptonight_monero, false, true>,
cryptonight_penta_hash<cryptonight_monero, true, true>,
+
cryptonight_double_hash<cryptonight_lite, false, false>,
cryptonight_double_hash<cryptonight_lite, true, false>,
cryptonight_double_hash<cryptonight_lite, false, true>,
@@ -550,6 +561,7 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
cryptonight_penta_hash<cryptonight_lite, true, false>,
cryptonight_penta_hash<cryptonight_lite, false, true>,
cryptonight_penta_hash<cryptonight_lite, true, true>,
+
cryptonight_double_hash<cryptonight, false, false>,
cryptonight_double_hash<cryptonight, true, false>,
cryptonight_double_hash<cryptonight, false, true>,
@@ -565,7 +577,41 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
cryptonight_penta_hash<cryptonight, false, false>,
cryptonight_penta_hash<cryptonight, true, false>,
cryptonight_penta_hash<cryptonight, false, true>,
- cryptonight_penta_hash<cryptonight, true, true>
+ cryptonight_penta_hash<cryptonight, true, true>,
+
+ cryptonight_double_hash<cryptonight_heavy, false, false>,
+ cryptonight_double_hash<cryptonight_heavy, true, false>,
+ cryptonight_double_hash<cryptonight_heavy, false, true>,
+ cryptonight_double_hash<cryptonight_heavy, true, true>,
+ cryptonight_triple_hash<cryptonight_heavy, false, false>,
+ cryptonight_triple_hash<cryptonight_heavy, true, false>,
+ cryptonight_triple_hash<cryptonight_heavy, false, true>,
+ cryptonight_triple_hash<cryptonight_heavy, true, true>,
+ cryptonight_quad_hash<cryptonight_heavy, false, false>,
+ cryptonight_quad_hash<cryptonight_heavy, true, false>,
+ cryptonight_quad_hash<cryptonight_heavy, false, true>,
+ cryptonight_quad_hash<cryptonight_heavy, true, true>,
+ cryptonight_penta_hash<cryptonight_heavy, false, false>,
+ cryptonight_penta_hash<cryptonight_heavy, true, false>,
+ cryptonight_penta_hash<cryptonight_heavy, false, true>,
+ cryptonight_penta_hash<cryptonight_heavy, true, true>,
+
+ cryptonight_double_hash<cryptonight_aeon, false, false>,
+ cryptonight_double_hash<cryptonight_aeon, true, false>,
+ cryptonight_double_hash<cryptonight_aeon, false, true>,
+ cryptonight_double_hash<cryptonight_aeon, true, true>,
+ cryptonight_triple_hash<cryptonight_aeon, false, false>,
+ cryptonight_triple_hash<cryptonight_aeon, true, false>,
+ cryptonight_triple_hash<cryptonight_aeon, false, true>,
+ cryptonight_triple_hash<cryptonight_aeon, true, true>,
+ cryptonight_quad_hash<cryptonight_aeon, false, false>,
+ cryptonight_quad_hash<cryptonight_aeon, true, false>,
+ cryptonight_quad_hash<cryptonight_aeon, false, true>,
+ cryptonight_quad_hash<cryptonight_aeon, true, true>,
+ cryptonight_penta_hash<cryptonight_aeon, false, false>,
+ cryptonight_penta_hash<cryptonight_aeon, true, false>,
+ cryptonight_penta_hash<cryptonight_aeon, false, true>,
+ cryptonight_penta_hash<cryptonight_aeon, true, true>
};
std::bitset<2> digit;
@@ -662,21 +708,15 @@ void minethd::multiway_work_main(cn_hash_fun_multi hash_fun_multi)
if(oWork.bNiceHash)
iNonce = *piNonce[0];
- if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero)
- {
- if(oWork.bWorkBlob[0] >= 7)
- hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight_monero);
- else
- hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight);
- }
+ auto miner_algo = ::jconf::inst()->GetMiningAlgo();
+ const bool time_to_fork =
+ ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && oWork.bWorkBlob[0] >= 7) ||
+ (miner_algo == cryptonight_heavy && oWork.bWorkBlob[0] >= 3);
- if(::jconf::inst()->GetMiningAlgo() == cryptonight_heavy)
- {
- if(oWork.bWorkBlob[0] >= 3)
- hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight_heavy);
- else
- hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight);
- }
+ if(time_to_fork)
+ hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
+ else
+ hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgoRoot());
while (globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo)
{
diff --git a/xmrstak/backend/cryptonight.hpp b/xmrstak/backend/cryptonight.hpp
index fe10a9f..538f1d8 100644
--- a/xmrstak/backend/cryptonight.hpp
+++ b/xmrstak/backend/cryptonight.hpp
@@ -9,7 +9,8 @@ enum xmrstak_algo
cryptonight = 1,
cryptonight_lite = 2,
cryptonight_monero = 3,
- cryptonight_heavy = 4
+ cryptonight_heavy = 4,
+ cryptonight_aeon = 5
};
// define aeon settings
@@ -40,6 +41,9 @@ inline constexpr size_t cn_select_memory<cryptonight_monero>() { return CRYPTONI
template<>
inline constexpr size_t cn_select_memory<cryptonight_heavy>() { return CRYPTONIGHT_HEAVY_MEMORY; }
+template<>
+inline constexpr size_t cn_select_memory<cryptonight_aeon>() { return CRYPTONIGHT_LITE_MEMORY; }
+
inline size_t cn_select_memory(xmrstak_algo algo)
{
@@ -53,6 +57,8 @@ inline size_t cn_select_memory(xmrstak_algo algo)
return CRYPTONIGHT_MEMORY;
case cryptonight_heavy:
return CRYPTONIGHT_HEAVY_MEMORY;
+ case cryptonight_aeon:
+ return CRYPTONIGHT_LITE_MEMORY;
default:
return 0;
}
@@ -73,6 +79,9 @@ inline constexpr uint32_t cn_select_mask<cryptonight_monero>() { return CRYPTONI
template<>
inline constexpr uint32_t cn_select_mask<cryptonight_heavy>() { return CRYPTONIGHT_HEAVY_MASK; }
+template<>
+inline constexpr uint32_t cn_select_mask<cryptonight_aeon>() { return CRYPTONIGHT_LITE_MASK; }
+
inline size_t cn_select_mask(xmrstak_algo algo)
{
switch(algo)
@@ -85,6 +94,8 @@ inline size_t cn_select_mask(xmrstak_algo algo)
return CRYPTONIGHT_MASK;
case cryptonight_heavy:
return CRYPTONIGHT_HEAVY_MASK;
+ case cryptonight_aeon:
+ return CRYPTONIGHT_LITE_MASK;
default:
return 0;
}
@@ -105,6 +116,9 @@ inline constexpr uint32_t cn_select_iter<cryptonight_monero>() { return CRYPTONI
template<>
inline constexpr uint32_t cn_select_iter<cryptonight_heavy>() { return CRYPTONIGHT_HEAVY_ITER; }
+template<>
+inline constexpr uint32_t cn_select_iter<cryptonight_aeon>() { return CRYPTONIGHT_LITE_ITER; }
+
inline size_t cn_select_iter(xmrstak_algo algo)
{
switch(algo)
@@ -117,6 +131,8 @@ inline size_t cn_select_iter(xmrstak_algo algo)
return CRYPTONIGHT_ITER;
case cryptonight_heavy:
return CRYPTONIGHT_HEAVY_ITER;
+ case cryptonight_aeon:
+ return CRYPTONIGHT_LITE_ITER;
default:
return 0;
}
diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp
index 153e4e3..f8e8409 100644
--- a/xmrstak/backend/nvidia/minethd.cpp
+++ b/xmrstak/backend/nvidia/minethd.cpp
@@ -238,14 +238,10 @@ void minethd::work_main()
cryptonight_ctx* cpu_ctx;
cpu_ctx = cpu::minethd::minethd_alloc_ctx();
auto miner_algo = ::jconf::inst()->GetMiningAlgo();
- cn_hash_fun hash_fun;
- if(miner_algo == cryptonight_monero || miner_algo == cryptonight_heavy)
- {
- // start with cryptonight and switch later if fork version is reached
- hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight);
- }
- else
- hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
+
+ // start with root algorithm and switch later if fork version is reached
+ cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, ::jconf::inst()->GetMiningAlgoRoot());
+
uint32_t iNonce;
globalStates::inst().iConsumeCnt++;
@@ -266,13 +262,12 @@ void minethd::work_main()
consume_work();
uint8_t new_version = oWork.getVersion();
- if(miner_algo == cryptonight_monero && version < 7 && new_version >= 7)
+ const bool time_to_fork =
+ ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version < 7 && new_version >= 7) ||
+ (miner_algo == cryptonight_heavy && version < 3 && new_version >= 3);
+ if(time_to_fork)
{
- hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_monero);
- }
- else if(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3)
- {
- hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy);
+ hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
}
version = new_version;
continue;
@@ -335,13 +330,12 @@ void minethd::work_main()
consume_work();
uint8_t new_version = oWork.getVersion();
- if(miner_algo == cryptonight_monero && version < 7 && new_version >= 7)
- {
- hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_monero);
- }
- else if(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3)
+ const bool time_to_fork =
+ ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version < 7 && new_version >= 7) ||
+ (miner_algo == cryptonight_heavy && version < 3 && new_version >= 3);
+ if(time_to_fork)
{
- hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy);
+ hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
}
version = new_version;
}
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
index ede578f..a1a9e15 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
@@ -231,7 +231,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
uint32_t t1[2], t2[2], res;
uint32_t tweak1_2[2];
- if (ALGO == cryptonight_monero)
+ if (ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
{
uint32_t * state = d_ctx_state + thread * 50;
tweak1_2[0] = (d_input[8] >> 24) | (d_input[9] << 8);
@@ -275,7 +275,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
t1[0] = shuffle<4>(sPtr,sub, d[x], 0);
const uint32_t z = d[0] ^ d[1];
- if(ALGO == cryptonight_monero)
+ if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
{
const uint32_t table = 0x75310U;
const uint32_t index = ((z >> 26) & 12) | ((z >> 23) & 2);
@@ -304,7 +304,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
res = *( (uint64_t *) t2 ) >> ( sub & 1 ? 32 : 0 );
- if(ALGO == cryptonight_monero)
+ if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon)
{
const uint32_t tweaked_res = tweak1_2[sub & 1] ^ res;
const uint32_t long_state_update = sub2 ? tweaked_res : res;
@@ -502,5 +502,12 @@ void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t
{
cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_lite>(ctx, startNonce);
}
+ else if(miner_algo == cryptonight_aeon)
+ {
+ if(version >= 7)
+ cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_aeon>(ctx, startNonce);
+ else
+ cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_lite>(ctx, startNonce);
+ }
}
diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp
index 713beb4..bbaeb8b 100644
--- a/xmrstak/jconf.cpp
+++ b/xmrstak/jconf.cpp
@@ -90,20 +90,21 @@ struct xmrstak_coin_algo
{
const char* coin_name;
xmrstak_algo algo;
+ xmrstak_algo algo_root;
const char* default_pool;
};
xmrstak_coin_algo coin_algos[] = {
- { "aeon", cryptonight_lite, "mine.aeon-pool.com:5555" },
- { "cryptonight", cryptonight, nullptr },
- { "cryptonight_lite", cryptonight_lite, nullptr },
- { "edollar", cryptonight, nullptr },
- { "electroneum", cryptonight, nullptr },
- { "graft", cryptonight, nullptr },
- { "intense", cryptonight, nullptr },
- { "karbo", cryptonight, nullptr },
- { "monero7", cryptonight_monero, "pool.usxmrpool.com:3333" },
- { "sumokoin", cryptonight_heavy, nullptr }
+ { "aeon7", cryptonight_aeon, cryptonight_lite, "mine.aeon-pool.com:5555" },
+ { "cryptonight", cryptonight, cryptonight, nullptr },
+ { "cryptonight_lite", cryptonight_lite, cryptonight_lite, nullptr },
+ { "edollar", cryptonight, cryptonight, nullptr },
+ { "electroneum", cryptonight, cryptonight, nullptr },
+ { "graft", cryptonight, cryptonight, nullptr },
+ { "intense", cryptonight, cryptonight, nullptr },
+ { "karbo", cryptonight, cryptonight, nullptr },
+ { "monero7", cryptonight_monero, cryptonight, "pool.usxmrpool.com:3333" },
+ { "sumokoin", cryptonight_heavy, cryptonight, nullptr }
};
constexpr size_t coin_alogo_size = (sizeof(coin_algos)/sizeof(coin_algos[0]));
@@ -630,6 +631,7 @@ bool jconf::parse_config(const char* sFilename, const char* sFilenamePools)
if(ctmp == coin_algos[i].coin_name)
{
mining_algo = coin_algos[i].algo;
+ mining_algo_root = coin_algos[i].algo_root;
break;
}
}
diff --git a/xmrstak/jconf.hpp b/xmrstak/jconf.hpp
index 6874d37..f81a3a2 100644
--- a/xmrstak/jconf.hpp
+++ b/xmrstak/jconf.hpp
@@ -49,6 +49,8 @@ public:
bool TlsSecureAlgos();
inline xmrstak_algo GetMiningAlgo() { return mining_algo; }
+
+ inline xmrstak_algo GetMiningAlgoRoot() { return mining_algo_root; }
std::string GetMiningCoin();
@@ -91,4 +93,5 @@ private:
bool bHaveAes;
xmrstak_algo mining_algo;
+ xmrstak_algo mining_algo_root;
};
diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp
index a620173..8520237 100644
--- a/xmrstak/misc/executor.cpp
+++ b/xmrstak/misc/executor.cpp
@@ -571,6 +571,13 @@ void executor::ex_main()
pools.emplace_front(0, "donate.xmr-stak.net:4444", "", "", "", 0.0, true, false, "", true);
break;
+ case cryptonight_aeon:
+ if(dev_tls)
+ pools.emplace_front(0, "donate.xmr-stak.net:7777", "", "", "", 0.0, true, true, "", true);
+ else
+ pools.emplace_front(0, "donate.xmr-stak.net:4444", "", "", "", 0.0, true, false, "", true);
+ break;
+
case cryptonight:
if(dev_tls)
pools.emplace_front(0, "donate.xmr-stak.net:6666", "", "", "", 0.0, true, true, "", false);
diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp
index 95bcc9c..5ea58a0 100644
--- a/xmrstak/net/jpsock.cpp
+++ b/xmrstak/net/jpsock.cpp
@@ -644,6 +644,9 @@ bool jpsock::cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bRes
case cryptonight_monero:
algo_name = "cryptonight-monero";
break;
+ case cryptonight_aeon:
+ algo_name = "cryptonight-aeon";
+ break;
default:
algo_name = "unknown";
break;
--
GitLab