From 81437eb07f72c203165b2ece3c7a77c6c665dc84 Mon Sep 17 00:00:00 2001 From: gnock Date: Tue, 5 Jun 2018 19:00:02 +0200 Subject: [PATCH 1/2] Add Cryptonight-fast --- xmrstak/backend/amd/amd_gpu/gpu.cpp | 2 +- .../backend/amd/amd_gpu/opencl/cryptonight.cl | 20 +++++------ .../backend/cpu/crypto/cryptonight_aesni.h | 32 +++++++++--------- xmrstak/backend/cpu/minethd.cpp | 33 ++++++++++++++++++- xmrstak/backend/cryptonight.hpp | 18 +++++++++- xmrstak/backend/nvidia/nvcc_code/cuda_core.cu | 12 ++++--- xmrstak/jconf.cpp | 3 +- xmrstak/net/jpsock.cpp | 3 ++ 8 files changed, 89 insertions(+), 34 deletions(-) diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index c664d531b..54bad53f8 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -1004,7 +1004,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 || miner_algo == cryptonight_aeon || miner_algo == cryptonight_ipbc || miner_algo == cryptonight_stellite) + if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon || miner_algo == cryptonight_ipbc || miner_algo == cryptonight_stellite || miner_algo == cryptonight_fast) { // Input if ((ret = clSetKernelArg(ctx->Kernels[kernel_storage][1], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS) diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl index c925c87a3..2c6529dd5 100644 --- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl +++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl @@ -553,8 +553,8 @@ __kernel void JOIN(cn0,ALGO)(__global ulong *input, __global uint4 *Scratchpad, __attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states, ulong Threads -// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite -#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7) +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_fast +#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8) , __global ulong *input #endif ) @@ -574,8 +574,8 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states } barrier(CLK_LOCAL_MEM_FENCE); -// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite -#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7) +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_fast +#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8) uint2 tweak1_2; #endif uint4 b_x; @@ -599,8 +599,8 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states b[1] = states[3] ^ states[7]; b_x = ((uint4 *)b)[0]; -// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite -#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7) +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_fast +#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8) tweak1_2 = as_uint2(input[4]); tweak1_2.s0 >>= 24; tweak1_2.s0 |= tweak1_2.s1 << 8; @@ -627,8 +627,8 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states ((uint4 *)c)[0] = AES_Round(AES0, AES1, AES2, AES3, ((uint4 *)c)[0], ((uint4 *)a)[0]); b_x ^= ((uint4 *)c)[0]; -// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite -#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7) +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_fast +#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8) uint table = 0x75310U; // cryptonight_stellite # if(ALGO == 7) @@ -646,8 +646,8 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states a[1] += c[0] * as_ulong2(tmp).s0; a[0] += mul_hi(c[0], as_ulong2(tmp).s0); -// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite -#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7) +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_fast +#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8) # if(ALGO == 6) uint2 ipbc_tmp = tweak1_2 ^ ((uint2 *)&(a[0]))[0]; diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h index e15c474a6..0ad9d15b7 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h +++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h @@ -432,7 +432,7 @@ inline void cryptonight_monero_tweak(uint64_t* mem_out, __m128i tmp) uint8_t x = static_cast(vh >> 24); static const uint16_t table = 0x7531; - if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_fast) { const uint8_t index = (((x >> 3) & 6) | (x & 1)) << 1; vh ^= ((table >> index) & 0x3) << 28; @@ -456,7 +456,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c constexpr size_t ITERATIONS = cn_select_iter(); constexpr size_t MEM = cn_select_memory(); - if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) && len < 43) { memset(output, 0, 32); return; @@ -465,7 +465,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) { monero_const = *reinterpret_cast(reinterpret_cast(input) + 35); monero_const ^= *(reinterpret_cast(ctx0->hash_state) + 24); @@ -494,7 +494,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) 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)); @@ -518,7 +518,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) { if(ALGO == cryptonight_ipbc) ((uint64_t*)&l0[idx0 & MASK])[1] = ah0 ^ monero_const ^ ((uint64_t*)&l0[idx0 & MASK])[0]; @@ -561,7 +561,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto constexpr size_t ITERATIONS = cn_select_iter(); constexpr size_t MEM = cn_select_memory(); - if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) && len < 43) { memset(output, 0, 64); return; @@ -571,7 +571,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) { monero_const_0 = *reinterpret_cast(reinterpret_cast(input) + 35); monero_const_0 ^= *(reinterpret_cast(ctx[0]->hash_state) + 24); @@ -609,7 +609,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) 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)); @@ -627,7 +627,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) 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)); @@ -648,7 +648,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) { if(ALGO == cryptonight_ipbc) ((uint64_t*)&l0[idx0 & MASK])[1] = axh0 ^ monero_const_0 ^ ((uint64_t*)&l0[idx0 & MASK])[0]; @@ -684,7 +684,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) { if(ALGO == cryptonight_ipbc) ((uint64_t*)&l1[idx1 & MASK])[1] = axh1 ^ monero_const_1 ^ ((uint64_t*)&l1[idx1 & MASK])[0]; @@ -736,7 +736,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) \ + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) \ cryptonight_monero_tweak((uint64_t*)ptr, b); \ else \ _mm_store_si128(ptr, b);\ @@ -751,7 +751,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) \ + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) \ { \ _mm_store_si128(ptr, _mm_xor_si128(a, mc)); \ if (ALGO == cryptonight_ipbc) \ @@ -782,7 +782,7 @@ void cryptonight_triple_hash(const void* input, size_t len, void* output, crypto constexpr size_t ITERATIONS = cn_select_iter(); constexpr size_t MEM = cn_select_memory(); - if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) && len < 43) { memset(output, 0, 32 * 3); return; @@ -876,7 +876,7 @@ void cryptonight_quad_hash(const void* input, size_t len, void* output, cryptoni constexpr size_t ITERATIONS = cn_select_iter(); constexpr size_t MEM = cn_select_memory(); - if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) && len < 43) { memset(output, 0, 32 * 4); return; @@ -985,7 +985,7 @@ void cryptonight_penta_hash(const void* input, size_t len, void* output, crypton constexpr size_t ITERATIONS = cn_select_iter(); constexpr size_t MEM = cn_select_memory(); - if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) && len < 43) { memset(output, 0, 32 * 5); return; diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp index 482c085e0..88617816e 100644 --- a/xmrstak/backend/cpu/minethd.cpp +++ b/xmrstak/backend/cpu/minethd.cpp @@ -292,6 +292,9 @@ bool minethd::self_test() else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_stellite) { } + else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_fast) + { + } for (int i = 0; i < MAX_N; i++) cryptonight_free_ctx(ctx[i]); @@ -377,6 +380,9 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr case cryptonight_stellite: algv = 6; break; + case cryptonight_fast: + algv = 7; + break; default: algv = 2; break; @@ -410,7 +416,11 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr cryptonight_hash, cryptonight_hash, cryptonight_hash, - cryptonight_hash + cryptonight_hash, + cryptonight_hash, + cryptonight_hash, + cryptonight_hash, + cryptonight_hash }; std::bitset<2> digit; @@ -555,6 +565,9 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, case cryptonight_stellite: algv = 6; break; + case cryptonight_fast: + algv = 7; + break; default: algv = 2; break; @@ -679,6 +692,24 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, cryptonight_penta_hash, cryptonight_penta_hash, cryptonight_penta_hash, + + + cryptonight_double_hash, + cryptonight_double_hash, + cryptonight_double_hash, + cryptonight_double_hash, + cryptonight_triple_hash, + cryptonight_triple_hash, + cryptonight_triple_hash, + cryptonight_triple_hash, + cryptonight_quad_hash, + cryptonight_quad_hash, + cryptonight_quad_hash, + cryptonight_quad_hash, + cryptonight_penta_hash, + cryptonight_penta_hash, + cryptonight_penta_hash, + cryptonight_penta_hash, }; std::bitset<2> digit; diff --git a/xmrstak/backend/cryptonight.hpp b/xmrstak/backend/cryptonight.hpp index 633ddf49b..b62e5170b 100644 --- a/xmrstak/backend/cryptonight.hpp +++ b/xmrstak/backend/cryptonight.hpp @@ -12,7 +12,8 @@ enum xmrstak_algo cryptonight_heavy = 4, cryptonight_aeon = 5, cryptonight_ipbc = 6, // equal to cryptonight_aeon with a small tweak in the miner code - cryptonight_stellite = 7 //equal to cryptonight_monero but with one tiny change + cryptonight_stellite = 7, //equal to cryptonight_monero but with one tiny change + cryptonight_fast = 8 //equal to cryptonight_monero but with less iterations, used by masari }; // define aeon settings @@ -28,6 +29,8 @@ constexpr size_t CRYPTONIGHT_HEAVY_MEMORY = 4 * 1024 * 1024; constexpr uint32_t CRYPTONIGHT_HEAVY_MASK = 0x3FFFF0; constexpr uint32_t CRYPTONIGHT_HEAVY_ITER = 0x40000; +constexpr uint32_t CRYPTONIGHT_FAST_ITER = 0x40000; + template inline constexpr size_t cn_select_memory() { return 0; } @@ -52,6 +55,9 @@ inline constexpr size_t cn_select_memory() { return CRYPTONIGH template<> inline constexpr size_t cn_select_memory() { return CRYPTONIGHT_MEMORY; } +template<> +inline constexpr size_t cn_select_memory() { return CRYPTONIGHT_MEMORY; } + inline size_t cn_select_memory(xmrstak_algo algo) { @@ -59,6 +65,7 @@ inline size_t cn_select_memory(xmrstak_algo algo) { case cryptonight_stellite: case cryptonight_monero: + case cryptonight_fast: case cryptonight: return CRYPTONIGHT_MEMORY; case cryptonight_ipbc: @@ -96,12 +103,16 @@ inline constexpr uint32_t cn_select_mask() { return CRYPTONIGH template<> inline constexpr uint32_t cn_select_mask() { return CRYPTONIGHT_MASK; } +template<> +inline constexpr uint32_t cn_select_mask() { return CRYPTONIGHT_MASK; } + inline size_t cn_select_mask(xmrstak_algo algo) { switch(algo) { case cryptonight_stellite: case cryptonight_monero: + case cryptonight_fast: case cryptonight: return CRYPTONIGHT_MASK; case cryptonight_ipbc: @@ -139,6 +150,9 @@ inline constexpr uint32_t cn_select_iter() { return CRYPTONIGH template<> inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_ITER; } +template<> +inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_FAST_ITER; } + inline size_t cn_select_iter(xmrstak_algo algo) { switch(algo) @@ -153,6 +167,8 @@ inline size_t cn_select_iter(xmrstak_algo algo) return CRYPTONIGHT_LITE_ITER; case cryptonight_heavy: return CRYPTONIGHT_HEAVY_ITER; + case cryptonight_fast: + return CRYPTONIGHT_FAST_ITER; default: return 0; } diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu index 57b6ad071..c9a888b07 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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) + if (ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) { uint32_t * state = d_ctx_state + thread * 50; tweak1_2[0] = (d_input[8] >> 24) | (d_input[9] << 8); @@ -275,10 +275,10 @@ __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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) { const uint32_t table = 0x75310U; - if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_fast) { const uint32_t index = ((z >> 26) & 12) | ((z >> 23) & 2); const uint32_t fork_7 = z ^ ((table >> index) & 0x30U) << 24; @@ -312,7 +312,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) { const uint32_t tweaked_res = tweak1_2[sub & 1] ^ res; uint32_t long_state_update = sub2 ? tweaked_res : res; @@ -515,5 +515,9 @@ void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t { cryptonight_core_gpu_hash(ctx, startNonce); } + else if(miner_algo == cryptonight_fast) + { + cryptonight_core_gpu_hash(ctx, startNonce); + } } diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp index b85ddd3cb..6ac3e0d99 100644 --- a/xmrstak/jconf.cpp +++ b/xmrstak/jconf.cpp @@ -92,6 +92,7 @@ xmrstak::coin_selection coins[] = { { "bbscoin", {cryptonight_monero, cryptonight, 3u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, { "croat", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, { "cryptonight", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "cryptonight_fast", {cryptonight_fast, cryptonight_fast, 0u}, {cryptonight_fast, cryptonight_fast, 0u}, nullptr }, { "cryptonight_heavy", {cryptonight_heavy, cryptonight_heavy, 0u}, {cryptonight_heavy, cryptonight_heavy, 0u}, nullptr }, { "cryptonight_lite", {cryptonight_aeon, cryptonight_lite, 255u}, {cryptonight_aeon, cryptonight_lite, 7u}, nullptr }, { "cryptonight_lite_v7", {cryptonight_lite, cryptonight_aeon, 255u}, {cryptonight_aeon, cryptonight_lite, 7u}, nullptr }, @@ -105,7 +106,7 @@ xmrstak::coin_selection coins[] = { { "intense", {cryptonight_monero, cryptonight, 4u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, { "ipbc", {cryptonight_aeon, cryptonight_ipbc, 255u}, {cryptonight_aeon, cryptonight_aeon, 255u}, nullptr }, { "karbo", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, - { "masari", {cryptonight_monero, cryptonight, 5u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "masari", {cryptonight_fast, cryptonight_monero, 7u}, {cryptonight_fast, cryptonight_fast, 0u}, nullptr }, { "monero7", {cryptonight_monero, cryptonight_monero, 0u}, {cryptonight_monero, cryptonight_monero, 0u}, "pool.usxmrpool.com:3333" }, { "stellite", {cryptonight_stellite, cryptonight_monero, 4u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, { "sumokoin", {cryptonight_heavy, cryptonight_heavy, 0u}, {cryptonight_heavy, cryptonight_heavy, 0u}, nullptr }, diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp index 6c41f2ba4..c2d2be73e 100644 --- a/xmrstak/net/jpsock.cpp +++ b/xmrstak/net/jpsock.cpp @@ -697,6 +697,9 @@ bool jpsock::cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bRes case cryptonight_heavy: algo_name = "cryptonight_heavy"; break; + case cryptonight_fast: + algo_name = "cryptonight_fast"; + break; default: algo_name = "unknown"; break; From adfbeb4c3c52b9f3f9465745766e6deb3ad11465 Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Thu, 7 Jun 2018 23:47:19 +0200 Subject: [PATCH 2/2] masari updates new algorithm - rename cryptonight_fast to cryptonight_masari - set dev pool to cryptonight_monero --- xmrstak/backend/amd/amd_gpu/gpu.cpp | 2 +- .../backend/amd/amd_gpu/opencl/cryptonight.cl | 10 ++-- .../backend/cpu/crypto/cryptonight_aesni.h | 32 ++++++------- xmrstak/backend/cpu/minethd.cpp | 46 +++++++++---------- xmrstak/backend/cryptonight.hpp | 18 ++++---- xmrstak/backend/nvidia/nvcc_code/cuda_core.cu | 12 ++--- xmrstak/jconf.cpp | 4 +- xmrstak/net/jpsock.cpp | 4 +- 8 files changed, 64 insertions(+), 64 deletions(-) diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index 54bad53f8..596ee2127 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -1004,7 +1004,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 || miner_algo == cryptonight_aeon || miner_algo == cryptonight_ipbc || miner_algo == cryptonight_stellite || miner_algo == cryptonight_fast) + if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon || miner_algo == cryptonight_ipbc || miner_algo == cryptonight_stellite || miner_algo == cryptonight_masari) { // Input if ((ret = clSetKernelArg(ctx->Kernels[kernel_storage][1], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS) diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl index 2c6529dd5..709a3659a 100644 --- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl +++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl @@ -553,7 +553,7 @@ __kernel void JOIN(cn0,ALGO)(__global ulong *input, __global uint4 *Scratchpad, __attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states, ulong Threads -// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_fast +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_masari #if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8) , __global ulong *input #endif @@ -574,7 +574,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states } barrier(CLK_LOCAL_MEM_FENCE); -// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_fast +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_masari #if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8) uint2 tweak1_2; #endif @@ -599,7 +599,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states b[1] = states[3] ^ states[7]; b_x = ((uint4 *)b)[0]; -// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_fast +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_masari #if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8) tweak1_2 = as_uint2(input[4]); tweak1_2.s0 >>= 24; @@ -627,7 +627,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states ((uint4 *)c)[0] = AES_Round(AES0, AES1, AES2, AES3, ((uint4 *)c)[0], ((uint4 *)a)[0]); b_x ^= ((uint4 *)c)[0]; -// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_fast +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_masari #if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8) uint table = 0x75310U; // cryptonight_stellite @@ -646,7 +646,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states a[1] += c[0] * as_ulong2(tmp).s0; a[0] += mul_hi(c[0], as_ulong2(tmp).s0); -// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_fast +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_masari #if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8) # if(ALGO == 6) diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h index 0ad9d15b7..364e32d1b 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h +++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h @@ -432,7 +432,7 @@ inline void cryptonight_monero_tweak(uint64_t* mem_out, __m128i tmp) uint8_t x = static_cast(vh >> 24); static const uint16_t table = 0x7531; - if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_fast) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_masari) { const uint8_t index = (((x >> 3) & 6) | (x & 1)) << 1; vh ^= ((table >> index) & 0x3) << 28; @@ -456,7 +456,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c constexpr size_t ITERATIONS = cn_select_iter(); constexpr size_t MEM = cn_select_memory(); - if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) && len < 43) { memset(output, 0, 32); return; @@ -465,7 +465,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) { monero_const = *reinterpret_cast(reinterpret_cast(input) + 35); monero_const ^= *(reinterpret_cast(ctx0->hash_state) + 24); @@ -494,7 +494,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) 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)); @@ -518,7 +518,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) { if(ALGO == cryptonight_ipbc) ((uint64_t*)&l0[idx0 & MASK])[1] = ah0 ^ monero_const ^ ((uint64_t*)&l0[idx0 & MASK])[0]; @@ -561,7 +561,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto constexpr size_t ITERATIONS = cn_select_iter(); constexpr size_t MEM = cn_select_memory(); - if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) && len < 43) { memset(output, 0, 64); return; @@ -571,7 +571,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) { monero_const_0 = *reinterpret_cast(reinterpret_cast(input) + 35); monero_const_0 ^= *(reinterpret_cast(ctx[0]->hash_state) + 24); @@ -609,7 +609,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) 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)); @@ -627,7 +627,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) 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)); @@ -648,7 +648,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) { if(ALGO == cryptonight_ipbc) ((uint64_t*)&l0[idx0 & MASK])[1] = axh0 ^ monero_const_0 ^ ((uint64_t*)&l0[idx0 & MASK])[0]; @@ -684,7 +684,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) { if(ALGO == cryptonight_ipbc) ((uint64_t*)&l1[idx1 & MASK])[1] = axh1 ^ monero_const_1 ^ ((uint64_t*)&l1[idx1 & MASK])[0]; @@ -736,7 +736,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) \ + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) \ cryptonight_monero_tweak((uint64_t*)ptr, b); \ else \ _mm_store_si128(ptr, b);\ @@ -751,7 +751,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) \ + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) \ { \ _mm_store_si128(ptr, _mm_xor_si128(a, mc)); \ if (ALGO == cryptonight_ipbc) \ @@ -782,7 +782,7 @@ void cryptonight_triple_hash(const void* input, size_t len, void* output, crypto constexpr size_t ITERATIONS = cn_select_iter(); constexpr size_t MEM = cn_select_memory(); - if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) && len < 43) { memset(output, 0, 32 * 3); return; @@ -876,7 +876,7 @@ void cryptonight_quad_hash(const void* input, size_t len, void* output, cryptoni constexpr size_t ITERATIONS = cn_select_iter(); constexpr size_t MEM = cn_select_memory(); - if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) && len < 43) { memset(output, 0, 32 * 4); return; @@ -985,7 +985,7 @@ void cryptonight_penta_hash(const void* input, size_t len, void* output, crypton constexpr size_t ITERATIONS = cn_select_iter(); constexpr size_t MEM = cn_select_memory(); - if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) && len < 43) { memset(output, 0, 32 * 5); return; diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp index 88617816e..b6a64fb48 100644 --- a/xmrstak/backend/cpu/minethd.cpp +++ b/xmrstak/backend/cpu/minethd.cpp @@ -292,7 +292,7 @@ bool minethd::self_test() else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_stellite) { } - else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_fast) + else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_masari) { } for (int i = 0; i < MAX_N; i++) @@ -380,7 +380,7 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr case cryptonight_stellite: algv = 6; break; - case cryptonight_fast: + case cryptonight_masari: algv = 7; break; default: @@ -417,10 +417,10 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr cryptonight_hash, cryptonight_hash, cryptonight_hash, - cryptonight_hash, - cryptonight_hash, - cryptonight_hash, - cryptonight_hash + cryptonight_hash, + cryptonight_hash, + cryptonight_hash, + cryptonight_hash }; std::bitset<2> digit; @@ -565,7 +565,7 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, case cryptonight_stellite: algv = 6; break; - case cryptonight_fast: + case cryptonight_masari: algv = 7; break; default: @@ -694,22 +694,22 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, cryptonight_penta_hash, - cryptonight_double_hash, - cryptonight_double_hash, - cryptonight_double_hash, - cryptonight_double_hash, - cryptonight_triple_hash, - cryptonight_triple_hash, - cryptonight_triple_hash, - cryptonight_triple_hash, - cryptonight_quad_hash, - cryptonight_quad_hash, - cryptonight_quad_hash, - cryptonight_quad_hash, - cryptonight_penta_hash, - cryptonight_penta_hash, - cryptonight_penta_hash, - cryptonight_penta_hash, + cryptonight_double_hash, + cryptonight_double_hash, + cryptonight_double_hash, + cryptonight_double_hash, + cryptonight_triple_hash, + cryptonight_triple_hash, + cryptonight_triple_hash, + cryptonight_triple_hash, + cryptonight_quad_hash, + cryptonight_quad_hash, + cryptonight_quad_hash, + cryptonight_quad_hash, + cryptonight_penta_hash, + cryptonight_penta_hash, + cryptonight_penta_hash, + cryptonight_penta_hash, }; std::bitset<2> digit; diff --git a/xmrstak/backend/cryptonight.hpp b/xmrstak/backend/cryptonight.hpp index b62e5170b..70655700d 100644 --- a/xmrstak/backend/cryptonight.hpp +++ b/xmrstak/backend/cryptonight.hpp @@ -13,7 +13,7 @@ enum xmrstak_algo cryptonight_aeon = 5, cryptonight_ipbc = 6, // equal to cryptonight_aeon with a small tweak in the miner code cryptonight_stellite = 7, //equal to cryptonight_monero but with one tiny change - cryptonight_fast = 8 //equal to cryptonight_monero but with less iterations, used by masari + cryptonight_masari = 8 //equal to cryptonight_monero but with less iterations, used by masari }; // define aeon settings @@ -29,7 +29,7 @@ constexpr size_t CRYPTONIGHT_HEAVY_MEMORY = 4 * 1024 * 1024; constexpr uint32_t CRYPTONIGHT_HEAVY_MASK = 0x3FFFF0; constexpr uint32_t CRYPTONIGHT_HEAVY_ITER = 0x40000; -constexpr uint32_t CRYPTONIGHT_FAST_ITER = 0x40000; +constexpr uint32_t CRYPTONIGHT_MASARI_ITER = 0x40000; template inline constexpr size_t cn_select_memory() { return 0; } @@ -56,7 +56,7 @@ template<> inline constexpr size_t cn_select_memory() { return CRYPTONIGHT_MEMORY; } template<> -inline constexpr size_t cn_select_memory() { return CRYPTONIGHT_MEMORY; } +inline constexpr size_t cn_select_memory() { return CRYPTONIGHT_MEMORY; } inline size_t cn_select_memory(xmrstak_algo algo) @@ -65,7 +65,7 @@ inline size_t cn_select_memory(xmrstak_algo algo) { case cryptonight_stellite: case cryptonight_monero: - case cryptonight_fast: + case cryptonight_masari: case cryptonight: return CRYPTONIGHT_MEMORY; case cryptonight_ipbc: @@ -104,7 +104,7 @@ template<> inline constexpr uint32_t cn_select_mask() { return CRYPTONIGHT_MASK; } template<> -inline constexpr uint32_t cn_select_mask() { return CRYPTONIGHT_MASK; } +inline constexpr uint32_t cn_select_mask() { return CRYPTONIGHT_MASK; } inline size_t cn_select_mask(xmrstak_algo algo) { @@ -112,7 +112,7 @@ inline size_t cn_select_mask(xmrstak_algo algo) { case cryptonight_stellite: case cryptonight_monero: - case cryptonight_fast: + case cryptonight_masari: case cryptonight: return CRYPTONIGHT_MASK; case cryptonight_ipbc: @@ -151,7 +151,7 @@ template<> inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_ITER; } template<> -inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_FAST_ITER; } +inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_MASARI_ITER; } inline size_t cn_select_iter(xmrstak_algo algo) { @@ -167,8 +167,8 @@ inline size_t cn_select_iter(xmrstak_algo algo) return CRYPTONIGHT_LITE_ITER; case cryptonight_heavy: return CRYPTONIGHT_HEAVY_ITER; - case cryptonight_fast: - return CRYPTONIGHT_FAST_ITER; + case cryptonight_masari: + return CRYPTONIGHT_MASARI_ITER; default: return 0; } diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu index c9a888b07..4d6d835b5 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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) + if (ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) { uint32_t * state = d_ctx_state + thread * 50; tweak1_2[0] = (d_input[8] >> 24) | (d_input[9] << 8); @@ -275,10 +275,10 @@ __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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) { const uint32_t table = 0x75310U; - if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_fast) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_masari) { const uint32_t index = ((z >> 26) & 12) | ((z >> 23) & 2); const uint32_t fork_7 = z ^ ((table >> index) & 0x30U) << 24; @@ -312,7 +312,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari) { const uint32_t tweaked_res = tweak1_2[sub & 1] ^ res; uint32_t long_state_update = sub2 ? tweaked_res : res; @@ -515,9 +515,9 @@ void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t { cryptonight_core_gpu_hash(ctx, startNonce); } - else if(miner_algo == cryptonight_fast) + else if(miner_algo == cryptonight_masari) { - cryptonight_core_gpu_hash(ctx, startNonce); + cryptonight_core_gpu_hash(ctx, startNonce); } } diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp index 6ac3e0d99..c49a1760a 100644 --- a/xmrstak/jconf.cpp +++ b/xmrstak/jconf.cpp @@ -92,7 +92,7 @@ xmrstak::coin_selection coins[] = { { "bbscoin", {cryptonight_monero, cryptonight, 3u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, { "croat", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, { "cryptonight", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, - { "cryptonight_fast", {cryptonight_fast, cryptonight_fast, 0u}, {cryptonight_fast, cryptonight_fast, 0u}, nullptr }, + { "cryptonight_masari", {cryptonight_monero, cryptonight_masari, 255u}, {cryptonight_monero, cryptonight_monero, 0u},nullptr }, { "cryptonight_heavy", {cryptonight_heavy, cryptonight_heavy, 0u}, {cryptonight_heavy, cryptonight_heavy, 0u}, nullptr }, { "cryptonight_lite", {cryptonight_aeon, cryptonight_lite, 255u}, {cryptonight_aeon, cryptonight_lite, 7u}, nullptr }, { "cryptonight_lite_v7", {cryptonight_lite, cryptonight_aeon, 255u}, {cryptonight_aeon, cryptonight_lite, 7u}, nullptr }, @@ -106,7 +106,7 @@ xmrstak::coin_selection coins[] = { { "intense", {cryptonight_monero, cryptonight, 4u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, { "ipbc", {cryptonight_aeon, cryptonight_ipbc, 255u}, {cryptonight_aeon, cryptonight_aeon, 255u}, nullptr }, { "karbo", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, - { "masari", {cryptonight_fast, cryptonight_monero, 7u}, {cryptonight_fast, cryptonight_fast, 0u}, nullptr }, + { "masari", {cryptonight_masari, cryptonight_monero, 7u}, {cryptonight_monero, cryptonight_monero, 0u},nullptr }, { "monero7", {cryptonight_monero, cryptonight_monero, 0u}, {cryptonight_monero, cryptonight_monero, 0u}, "pool.usxmrpool.com:3333" }, { "stellite", {cryptonight_stellite, cryptonight_monero, 4u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, { "sumokoin", {cryptonight_heavy, cryptonight_heavy, 0u}, {cryptonight_heavy, cryptonight_heavy, 0u}, nullptr }, diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp index c2d2be73e..cf04c8605 100644 --- a/xmrstak/net/jpsock.cpp +++ b/xmrstak/net/jpsock.cpp @@ -697,8 +697,8 @@ bool jpsock::cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bRes case cryptonight_heavy: algo_name = "cryptonight_heavy"; break; - case cryptonight_fast: - algo_name = "cryptonight_fast"; + case cryptonight_masari: + algo_name = "cryptonight_masari"; break; default: algo_name = "unknown";