diff options
Diffstat (limited to 'xmrstak/backend')
| -rw-r--r-- | xmrstak/backend/amd/amd_gpu/gpu.cpp | 6 | ||||
| -rw-r--r-- | xmrstak/backend/amd/minethd.cpp | 33 | ||||
| -rw-r--r-- | xmrstak/backend/cpu/crypto/cryptonight_aesni.h | 30 | ||||
| -rw-r--r-- | xmrstak/backend/cpu/minethd.cpp | 100 | ||||
| -rw-r--r-- | xmrstak/backend/cryptonight.hpp | 18 | ||||
| -rw-r--r-- | xmrstak/backend/nvidia/minethd.cpp | 34 | ||||
| -rw-r--r-- | xmrstak/backend/nvidia/nvcc_code/cuda_core.cu | 13 |
7 files changed, 142 insertions, 92 deletions
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index 79befb7..4a82fab 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -940,7 +940,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar /// @todo only activate if currency is monero int cn_kernel_offset = 0; - if(miner_algo == cryptonight_monero && version >= 7) + if((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon ) && version >= 7) { cn_kernel_offset = 6; } @@ -966,7 +966,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar return(ERR_OCL_API); } - if(miner_algo == cryptonight_monero && version >= 7) + if((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon ) && version >= 7) { // Input if ((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS) @@ -1134,7 +1134,7 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo, size_t tmpNonce = ctx->Nonce; /// @todo only activate if currency is monero int cn_kernel_offset = 0; - if(miner_algo == cryptonight_monero && version >= 7) + if((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version >= 7) { cn_kernel_offset = 6; } diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp index d461aad..775aa08 100644 --- a/xmrstak/backend/amd/minethd.cpp +++ b/xmrstak/backend/amd/minethd.cpp @@ -192,14 +192,9 @@ void minethd::work_main() cryptonight_ctx* cpu_ctx; cpu_ctx = cpu::minethd::minethd_alloc_ctx(); auto miner_algo = ::jconf::inst()->GetMiningAlgo(); - cn_hash_fun hash_fun; - if(miner_algo == cryptonight_monero || miner_algo == cryptonight_heavy) - { - // start with cryptonight and switch later if fork version is reached - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight); - } - else - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); + + // start with root algorithm and switch later if fork version is reached + cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, ::jconf::inst()->GetMiningAlgoRoot()); globalStates::inst().iConsumeCnt++; @@ -219,13 +214,12 @@ void minethd::work_main() consume_work(); uint8_t new_version = oWork.getVersion(); - if(miner_algo == cryptonight_monero && version < 7 && new_version >= 7) - { - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_monero); - } - else if(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3) + const bool time_to_fork = + ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version < 7 && new_version >= 7) || + (miner_algo == cryptonight_heavy && version < 3 && new_version >= 3); + if(time_to_fork) { - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy); + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); } version = new_version; continue; @@ -281,13 +275,12 @@ void minethd::work_main() consume_work(); uint8_t new_version = oWork.getVersion(); - if(miner_algo == cryptonight_monero && version < 7 && new_version >= 7) - { - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_monero); - } - else if(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3) + const bool time_to_fork = + ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version < 7 && new_version >= 7) || + (miner_algo == cryptonight_heavy && version < 3 && new_version >= 3); + if(time_to_fork) { - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy); + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); } version = new_version; } diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h index 85373e8..5203de8 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h +++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h @@ -444,7 +444,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c constexpr size_t ITERATIONS = cn_select_iter<ALGO>(); constexpr size_t MEM = cn_select_memory<ALGO>(); - if(ALGO == cryptonight_monero && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43) { memset(output, 0, 32); return; @@ -453,7 +453,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c keccak((const uint8_t *)input, len, ctx0->hash_state, 200); uint64_t monero_const; - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) { monero_const = *reinterpret_cast<const uint64_t*>(reinterpret_cast<const uint8_t*>(input) + 35); monero_const ^= *(reinterpret_cast<const uint64_t*>(ctx0->hash_state) + 24); @@ -482,7 +482,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c else cx = _mm_aesenc_si128(cx, _mm_set_epi64x(ah0, al0)); - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) cryptonight_monero_tweak((uint64_t*)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx)); else _mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx)); @@ -506,7 +506,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c _mm_prefetch((const char*)&l0[al0 & MASK], _MM_HINT_T0); ah0 += lo; - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) ((uint64_t*)&l0[idx0 & MASK])[1] = ah0 ^ monero_const; else ((uint64_t*)&l0[idx0 & MASK])[1] = ah0; @@ -544,7 +544,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto constexpr size_t ITERATIONS = cn_select_iter<ALGO>(); constexpr size_t MEM = cn_select_memory<ALGO>(); - if(ALGO == cryptonight_monero && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43) { memset(output, 0, 64); return; @@ -554,7 +554,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto keccak((const uint8_t *)input+len, len, ctx[1]->hash_state, 200); uint64_t monero_const_0, monero_const_1; - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) { monero_const_0 = *reinterpret_cast<const uint64_t*>(reinterpret_cast<const uint8_t*>(input) + 35); monero_const_0 ^= *(reinterpret_cast<const uint64_t*>(ctx[0]->hash_state) + 24); @@ -592,7 +592,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto else cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh0, axl0)); - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) cryptonight_monero_tweak((uint64_t*)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx)); else _mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx)); @@ -610,7 +610,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto else cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh1, axl1)); - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) cryptonight_monero_tweak((uint64_t*)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx)); else _mm_store_si128((__m128i *)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx)); @@ -631,7 +631,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto axh0 += lo; ((uint64_t*)&l0[idx0 & MASK])[0] = axl0; - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) ((uint64_t*)&l0[idx0 & MASK])[1] = axh0 ^ monero_const_0; else ((uint64_t*)&l0[idx0 & MASK])[1] = axh0; @@ -662,7 +662,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto axh1 += lo; ((uint64_t*)&l1[idx1 & MASK])[0] = axl1; - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) ((uint64_t*)&l1[idx1 & MASK])[1] = axh1 ^ monero_const_1; else ((uint64_t*)&l1[idx1 & MASK])[1] = axh1; @@ -709,7 +709,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto else \ c = _mm_aesenc_si128(c, a); \ b = _mm_xor_si128(b, c); \ - if(ALGO == cryptonight_monero) \ + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) \ cryptonight_monero_tweak((uint64_t*)ptr, b); \ else \ _mm_store_si128(ptr, b);\ @@ -724,7 +724,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto #define CN_STEP4(a, b, c, l, mc, ptr, idx) \ lo = _umul128(idx, _mm_cvtsi128_si64(b), &hi); \ a = _mm_add_epi64(a, _mm_set_epi64x(lo, hi)); \ - if(ALGO == cryptonight_monero) \ + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) \ _mm_store_si128(ptr, _mm_xor_si128(a, mc)); \ else \ _mm_store_si128(ptr, a);\ @@ -751,7 +751,7 @@ void cryptonight_triple_hash(const void* input, size_t len, void* output, crypto constexpr size_t ITERATIONS = cn_select_iter<ALGO>(); constexpr size_t MEM = cn_select_memory<ALGO>(); - if(ALGO == cryptonight_monero && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43) { memset(output, 0, 32 * 3); return; @@ -845,7 +845,7 @@ void cryptonight_quad_hash(const void* input, size_t len, void* output, cryptoni constexpr size_t ITERATIONS = cn_select_iter<ALGO>(); constexpr size_t MEM = cn_select_memory<ALGO>(); - if(ALGO == cryptonight_monero && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43) { memset(output, 0, 32 * 4); return; @@ -954,7 +954,7 @@ void cryptonight_penta_hash(const void* input, size_t len, void* output, crypton constexpr size_t ITERATIONS = cn_select_iter<ALGO>(); constexpr size_t MEM = cn_select_memory<ALGO>(); - if(ALGO == cryptonight_monero && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43) { memset(output, 0, 32 * 5); return; diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp index d5324f5..b25d91e 100644 --- a/xmrstak/backend/cpu/minethd.cpp +++ b/xmrstak/backend/cpu/minethd.cpp @@ -282,6 +282,9 @@ bool minethd::self_test() else if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero) { } + else if(::jconf::inst()->GetMiningAlgo() == cryptonight_aeon) + { + } for (int i = 0; i < MAX_N; i++) cryptonight_free_ctx(ctx[i]); @@ -362,6 +365,9 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr case cryptonight_heavy: algv = 3; break; + case cryptonight_aeon: + algv = 4; + break; default: algv = 2; break; @@ -383,7 +389,11 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr cryptonight_hash<cryptonight_heavy, false, false>, cryptonight_hash<cryptonight_heavy, true, false>, cryptonight_hash<cryptonight_heavy, false, true>, - cryptonight_hash<cryptonight_heavy, true, true> + cryptonight_hash<cryptonight_heavy, true, true>, + cryptonight_hash<cryptonight_aeon, false, false>, + cryptonight_hash<cryptonight_aeon, true, false>, + cryptonight_hash<cryptonight_aeon, false, true>, + cryptonight_hash<cryptonight_aeon, true, true> }; std::bitset<2> digit; @@ -443,21 +453,15 @@ void minethd::work_main() if(oWork.bNiceHash) result.iNonce = *piNonce; - if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero) - { - if(oWork.bWorkBlob[0] >= 7) - hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight_monero); - else - hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight); - } + auto miner_algo = ::jconf::inst()->GetMiningAlgo(); + const bool time_to_fork = + ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && oWork.bWorkBlob[0] >= 7) || + (miner_algo == cryptonight_heavy && oWork.bWorkBlob[0] >= 3); - if(::jconf::inst()->GetMiningAlgo() == cryptonight_heavy) - { - if(oWork.bWorkBlob[0] >= 3) - hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight_heavy); - else - hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight); - } + if(time_to_fork) + hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo); + else + hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgoRoot()); while(globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo) { @@ -508,6 +512,12 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, case cryptonight_monero: algv = 0; break; + case cryptonight_heavy: + algv = 3; + break; + case cryptonight_aeon: + algv = 4; + break; default: algv = 2; break; @@ -530,6 +540,7 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, cryptonight_penta_hash<cryptonight_monero, true, false>, cryptonight_penta_hash<cryptonight_monero, false, true>, cryptonight_penta_hash<cryptonight_monero, true, true>, + cryptonight_double_hash<cryptonight_lite, false, false>, cryptonight_double_hash<cryptonight_lite, true, false>, cryptonight_double_hash<cryptonight_lite, false, true>, @@ -546,6 +557,7 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, cryptonight_penta_hash<cryptonight_lite, true, false>, cryptonight_penta_hash<cryptonight_lite, false, true>, cryptonight_penta_hash<cryptonight_lite, true, true>, + cryptonight_double_hash<cryptonight, false, false>, cryptonight_double_hash<cryptonight, true, false>, cryptonight_double_hash<cryptonight, false, true>, @@ -561,7 +573,41 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, cryptonight_penta_hash<cryptonight, false, false>, cryptonight_penta_hash<cryptonight, true, false>, cryptonight_penta_hash<cryptonight, false, true>, - cryptonight_penta_hash<cryptonight, true, true> + cryptonight_penta_hash<cryptonight, true, true>, + + cryptonight_double_hash<cryptonight_heavy, false, false>, + cryptonight_double_hash<cryptonight_heavy, true, false>, + cryptonight_double_hash<cryptonight_heavy, false, true>, + cryptonight_double_hash<cryptonight_heavy, true, true>, + cryptonight_triple_hash<cryptonight_heavy, false, false>, + cryptonight_triple_hash<cryptonight_heavy, true, false>, + cryptonight_triple_hash<cryptonight_heavy, false, true>, + cryptonight_triple_hash<cryptonight_heavy, true, true>, + cryptonight_quad_hash<cryptonight_heavy, false, false>, + cryptonight_quad_hash<cryptonight_heavy, true, false>, + cryptonight_quad_hash<cryptonight_heavy, false, true>, + cryptonight_quad_hash<cryptonight_heavy, true, true>, + cryptonight_penta_hash<cryptonight_heavy, false, false>, + cryptonight_penta_hash<cryptonight_heavy, true, false>, + cryptonight_penta_hash<cryptonight_heavy, false, true>, + cryptonight_penta_hash<cryptonight_heavy, true, true>, + + cryptonight_double_hash<cryptonight_aeon, false, false>, + cryptonight_double_hash<cryptonight_aeon, true, false>, + cryptonight_double_hash<cryptonight_aeon, false, true>, + cryptonight_double_hash<cryptonight_aeon, true, true>, + cryptonight_triple_hash<cryptonight_aeon, false, false>, + cryptonight_triple_hash<cryptonight_aeon, true, false>, + cryptonight_triple_hash<cryptonight_aeon, false, true>, + cryptonight_triple_hash<cryptonight_aeon, true, true>, + cryptonight_quad_hash<cryptonight_aeon, false, false>, + cryptonight_quad_hash<cryptonight_aeon, true, false>, + cryptonight_quad_hash<cryptonight_aeon, false, true>, + cryptonight_quad_hash<cryptonight_aeon, true, true>, + cryptonight_penta_hash<cryptonight_aeon, false, false>, + cryptonight_penta_hash<cryptonight_aeon, true, false>, + cryptonight_penta_hash<cryptonight_aeon, false, true>, + cryptonight_penta_hash<cryptonight_aeon, true, true> }; std::bitset<2> digit; @@ -658,21 +704,15 @@ void minethd::multiway_work_main(cn_hash_fun_multi hash_fun_multi) if(oWork.bNiceHash) iNonce = *piNonce[0]; - if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero) - { - if(oWork.bWorkBlob[0] >= 7) - hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight_monero); - else - hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight); - } + auto miner_algo = ::jconf::inst()->GetMiningAlgo(); + const bool time_to_fork = + ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && oWork.bWorkBlob[0] >= 7) || + (miner_algo == cryptonight_heavy && oWork.bWorkBlob[0] >= 3); - if(::jconf::inst()->GetMiningAlgo() == cryptonight_heavy) - { - if(oWork.bWorkBlob[0] >= 3) - hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight_heavy); - else - hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight); - } + if(time_to_fork) + hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo); + else + hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgoRoot()); while (globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo) { diff --git a/xmrstak/backend/cryptonight.hpp b/xmrstak/backend/cryptonight.hpp index fe10a9f..538f1d8 100644 --- a/xmrstak/backend/cryptonight.hpp +++ b/xmrstak/backend/cryptonight.hpp @@ -9,7 +9,8 @@ enum xmrstak_algo cryptonight = 1, cryptonight_lite = 2, cryptonight_monero = 3, - cryptonight_heavy = 4 + cryptonight_heavy = 4, + cryptonight_aeon = 5 }; // define aeon settings @@ -40,6 +41,9 @@ inline constexpr size_t cn_select_memory<cryptonight_monero>() { return CRYPTONI template<> inline constexpr size_t cn_select_memory<cryptonight_heavy>() { return CRYPTONIGHT_HEAVY_MEMORY; } +template<> +inline constexpr size_t cn_select_memory<cryptonight_aeon>() { return CRYPTONIGHT_LITE_MEMORY; } + inline size_t cn_select_memory(xmrstak_algo algo) { @@ -53,6 +57,8 @@ inline size_t cn_select_memory(xmrstak_algo algo) return CRYPTONIGHT_MEMORY; case cryptonight_heavy: return CRYPTONIGHT_HEAVY_MEMORY; + case cryptonight_aeon: + return CRYPTONIGHT_LITE_MEMORY; default: return 0; } @@ -73,6 +79,9 @@ inline constexpr uint32_t cn_select_mask<cryptonight_monero>() { return CRYPTONI template<> inline constexpr uint32_t cn_select_mask<cryptonight_heavy>() { return CRYPTONIGHT_HEAVY_MASK; } +template<> +inline constexpr uint32_t cn_select_mask<cryptonight_aeon>() { return CRYPTONIGHT_LITE_MASK; } + inline size_t cn_select_mask(xmrstak_algo algo) { switch(algo) @@ -85,6 +94,8 @@ inline size_t cn_select_mask(xmrstak_algo algo) return CRYPTONIGHT_MASK; case cryptonight_heavy: return CRYPTONIGHT_HEAVY_MASK; + case cryptonight_aeon: + return CRYPTONIGHT_LITE_MASK; default: return 0; } @@ -105,6 +116,9 @@ inline constexpr uint32_t cn_select_iter<cryptonight_monero>() { return CRYPTONI template<> inline constexpr uint32_t cn_select_iter<cryptonight_heavy>() { return CRYPTONIGHT_HEAVY_ITER; } +template<> +inline constexpr uint32_t cn_select_iter<cryptonight_aeon>() { return CRYPTONIGHT_LITE_ITER; } + inline size_t cn_select_iter(xmrstak_algo algo) { switch(algo) @@ -117,6 +131,8 @@ inline size_t cn_select_iter(xmrstak_algo algo) return CRYPTONIGHT_ITER; case cryptonight_heavy: return CRYPTONIGHT_HEAVY_ITER; + case cryptonight_aeon: + return CRYPTONIGHT_LITE_ITER; default: return 0; } diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp index 75f5fb1..79297b0 100644 --- a/xmrstak/backend/nvidia/minethd.cpp +++ b/xmrstak/backend/nvidia/minethd.cpp @@ -238,14 +238,10 @@ void minethd::work_main() cryptonight_ctx* cpu_ctx; cpu_ctx = cpu::minethd::minethd_alloc_ctx(); auto miner_algo = ::jconf::inst()->GetMiningAlgo(); - cn_hash_fun hash_fun; - if(miner_algo == cryptonight_monero || miner_algo == cryptonight_heavy) - { - // start with cryptonight and switch later if fork version is reached - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight); - } - else - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); + + // start with root algorithm and switch later if fork version is reached + cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, ::jconf::inst()->GetMiningAlgoRoot()); + uint32_t iNonce; globalStates::inst().iConsumeCnt++; @@ -266,13 +262,12 @@ void minethd::work_main() consume_work(); uint8_t new_version = oWork.getVersion(); - if(miner_algo == cryptonight_monero && version < 7 && new_version >= 7) + const bool time_to_fork = + ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version < 7 && new_version >= 7) || + (miner_algo == cryptonight_heavy && version < 3 && new_version >= 3); + if(time_to_fork) { - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_monero); - } - else if(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3) - { - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy); + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); } version = new_version; continue; @@ -335,13 +330,12 @@ void minethd::work_main() consume_work(); uint8_t new_version = oWork.getVersion(); - if(miner_algo == cryptonight_monero && version < 7 && new_version >= 7) - { - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_monero); - } - else if(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3) + const bool time_to_fork = + ((miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) && version < 7 && new_version >= 7) || + (miner_algo == cryptonight_heavy && version < 3 && new_version >= 3); + if(time_to_fork) { - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy); + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); } version = new_version; } diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu index ede578f..a1a9e15 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu @@ -231,7 +231,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti uint32_t t1[2], t2[2], res; uint32_t tweak1_2[2]; - if (ALGO == cryptonight_monero) + if (ALGO == cryptonight_monero || ALGO == cryptonight_aeon) { uint32_t * state = d_ctx_state + thread * 50; tweak1_2[0] = (d_input[8] >> 24) | (d_input[9] << 8); @@ -275,7 +275,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti t1[0] = shuffle<4>(sPtr,sub, d[x], 0); const uint32_t z = d[0] ^ d[1]; - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) { const uint32_t table = 0x75310U; const uint32_t index = ((z >> 26) & 12) | ((z >> 23) & 2); @@ -304,7 +304,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti res = *( (uint64_t *) t2 ) >> ( sub & 1 ? 32 : 0 ); - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) { const uint32_t tweaked_res = tweak1_2[sub & 1] ^ res; const uint32_t long_state_update = sub2 ? tweaked_res : res; @@ -502,5 +502,12 @@ void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t { cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_lite>(ctx, startNonce); } + else if(miner_algo == cryptonight_aeon) + { + if(version >= 7) + cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_aeon>(ctx, startNonce); + else + cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_lite>(ctx, startNonce); + } } |
