From 84471e3d48b0423410384d41c2ce467b19f2b73a Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Fri, 27 Apr 2018 21:45:40 +0200 Subject: support stellite v4 fork solve #1494 - add algorithm `cryptonight_v7_stellite` (internal named: `cryptonight_stellite`) --- xmrstak/backend/amd/amd_gpu/gpu.cpp | 2 +- xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl | 22 ++++++--- xmrstak/backend/cpu/crypto/cryptonight_aesni.h | 56 ++++++++++++++--------- xmrstak/backend/cpu/minethd.cpp | 34 +++++++++++++- xmrstak/backend/cryptonight.hpp | 15 +++++- xmrstak/backend/nvidia/nvcc_code/cuda_core.cu | 25 +++++++--- xmrstak/jconf.cpp | 3 +- xmrstak/misc/executor.cpp | 2 +- xmrstak/net/jpsock.cpp | 3 ++ 9 files changed, 121 insertions(+), 41 deletions(-) (limited to 'xmrstak') diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index 736b6ae..59f654f 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) + if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon || miner_algo == cryptonight_ipbc || cryptonight_stellite) { // 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 0738c04..9e2f03c 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 -#if(ALGO == 3 || ALGO == 5 || ALGO == 6) +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite +#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7) , __global ulong *input #endif ) @@ -574,7 +574,8 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states } barrier(CLK_LOCAL_MEM_FENCE); -#if(ALGO == 3 || ALGO == 5 || ALGO == 6) +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite +#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7) uint2 tweak1_2; #endif uint4 b_x; @@ -598,7 +599,8 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states b[1] = states[3] ^ states[7]; b_x = ((uint4 *)b)[0]; -#if(ALGO == 3 || ALGO == 5 || ALGO == 6) +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite +#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7) tweak1_2 = as_uint2(input[4]); tweak1_2.s0 >>= 24; tweak1_2.s0 |= tweak1_2.s1 << 8; @@ -625,9 +627,15 @@ __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]; -#if(ALGO == 3 || ALGO == 5 || ALGO == 6) +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite +#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7) uint table = 0x75310U; +// cryptonight_stellite +# if(ALGO == 7) + uint index = ((b_x.s2 >> 27) & 12) | ((b_x.s2 >> 23) & 2); +# else uint index = ((b_x.s2 >> 26) & 12) | ((b_x.s2 >> 23) & 2); +# endif b_x.s2 ^= ((table >> index) & 0x30U) << 24; #endif Scratchpad[IDX((idx0 & MASK) >> 4)] = b_x; @@ -638,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); - -#if(ALGO == 3 || ALGO == 5 || ALGO == 6) +// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite +#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7) # 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 e2f4d88..da0a8af 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h +++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h @@ -422,6 +422,7 @@ void cn_implode_scratchpad(const __m128i* input, __m128i* output) _mm_store_si128(output + 11, xout7); } +template inline void cryptonight_monero_tweak(uint64_t* mem_out, __m128i tmp) { mem_out[0] = _mm_cvtsi128_si64(tmp); @@ -431,10 +432,21 @@ inline void cryptonight_monero_tweak(uint64_t* mem_out, __m128i tmp) uint8_t x = static_cast(vh >> 24); static const uint16_t table = 0x7531; - const uint8_t index = (((x >> 3) & 6) | (x & 1)) << 1; - vh ^= ((table >> index) & 0x3) << 28; + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) + { + const uint8_t index = (((x >> 3) & 6) | (x & 1)) << 1; + vh ^= ((table >> index) & 0x3) << 28; + + mem_out[1] = vh; + } + else if(ALGO == cryptonight_stellite) + { + const uint8_t index = (((x >> 4) & 6) | (x & 1)) << 1; + vh ^= ((table >> index) & 0x3) << 28; + + mem_out[1] = vh; + } - mem_out[1] = vh; } template @@ -444,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) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) && len < 43) { memset(output, 0, 32); return; @@ -453,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) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) { monero_const = *reinterpret_cast(reinterpret_cast(input) + 35); monero_const ^= *(reinterpret_cast(ctx0->hash_state) + 24); @@ -482,8 +494,8 @@ 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) - cryptonight_monero_tweak((uint64_t*)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx)); + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) + 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 +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) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) { if(ALGO == cryptonight_ipbc) ((uint64_t*)&l0[idx0 & MASK])[1] = ah0 ^ monero_const ^ ((uint64_t*)&l0[idx0 & MASK])[0]; @@ -549,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) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) && len < 43) { memset(output, 0, 64); return; @@ -559,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) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) { monero_const_0 = *reinterpret_cast(reinterpret_cast(input) + 35); monero_const_0 ^= *(reinterpret_cast(ctx[0]->hash_state) + 24); @@ -597,8 +609,8 @@ 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) - cryptonight_monero_tweak((uint64_t*)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx)); + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) + 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)); @@ -615,8 +627,8 @@ 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) - cryptonight_monero_tweak((uint64_t*)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx)); + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) + 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)); @@ -636,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) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) { if(ALGO == cryptonight_ipbc) ((uint64_t*)&l0[idx0 & MASK])[1] = axh0 ^ monero_const_0 ^ ((uint64_t*)&l0[idx0 & MASK])[0]; @@ -672,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) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) { if(ALGO == cryptonight_ipbc) ((uint64_t*)&l1[idx1 & MASK])[1] = axh1 ^ monero_const_1 ^ ((uint64_t*)&l1[idx1 & MASK])[0]; @@ -724,8 +736,8 @@ 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) \ - cryptonight_monero_tweak((uint64_t*)ptr, b); \ + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) \ + cryptonight_monero_tweak((uint64_t*)ptr, b); \ else \ _mm_store_si128(ptr, b);\ @@ -739,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) \ + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) \ { \ _mm_store_si128(ptr, _mm_xor_si128(a, mc)); \ if (ALGO == cryptonight_ipbc) \ @@ -770,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) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) && len < 43) { memset(output, 0, 32 * 3); return; @@ -864,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) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) && len < 43) { memset(output, 0, 32 * 4); return; @@ -973,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) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) && len < 43) { memset(output, 0, 32 * 5); return; diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp index bbb4eec..5753fd6 100644 --- a/xmrstak/backend/cpu/minethd.cpp +++ b/xmrstak/backend/cpu/minethd.cpp @@ -289,6 +289,9 @@ bool minethd::self_test() else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_ipbc) { } + else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_stellite) + { + } for (int i = 0; i < MAX_N; i++) cryptonight_free_ctx(ctx[i]); @@ -374,6 +377,9 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr case cryptonight_ipbc: algv = 5; break; + case cryptonight_stellite: + algv = 6; + break; default: algv = 2; break; @@ -403,7 +409,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; @@ -543,6 +553,9 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, case cryptonight_ipbc: algv = 5; break; + case cryptonight_stellite: + algv = 6; + break; default: algv = 2; break; @@ -649,7 +662,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_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 b175cda..c0ad0cd 100644 --- a/xmrstak/backend/cryptonight.hpp +++ b/xmrstak/backend/cryptonight.hpp @@ -11,7 +11,8 @@ enum xmrstak_algo cryptonight_monero = 3, cryptonight_heavy = 4, cryptonight_aeon = 5, - cryptonight_ipbc = 6 // equal to cryptonight_aeon with a small tweak in the miner code + 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 }; // define aeon settings @@ -48,11 +49,15 @@ inline constexpr size_t cn_select_memory() { return CRYPTONIGH template<> inline constexpr size_t cn_select_memory() { return CRYPTONIGHT_LITE_MEMORY; } +template<> +inline constexpr size_t cn_select_memory() { return CRYPTONIGHT_MEMORY; } + inline size_t cn_select_memory(xmrstak_algo algo) { switch(algo) { + case cryptonight_stellite: case cryptonight_monero: case cryptonight: return CRYPTONIGHT_MEMORY; @@ -88,10 +93,14 @@ inline constexpr uint32_t cn_select_mask() { return CRYPTONIGH template<> inline constexpr uint32_t cn_select_mask() { return CRYPTONIGHT_LITE_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: return CRYPTONIGHT_MASK; @@ -127,10 +136,14 @@ inline constexpr uint32_t cn_select_iter() { return CRYPTONIGH template<> inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_LITE_ITER; } +template<> +inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_ITER; } + inline size_t cn_select_iter(xmrstak_algo algo) { switch(algo) { + case cryptonight_stellite: case cryptonight_monero: case cryptonight: return CRYPTONIGHT_ITER; diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu index 3bb910e..4eacfb6 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) + if (ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) { uint32_t * state = d_ctx_state + thread * 50; tweak1_2[0] = (d_input[8] >> 24) | (d_input[9] << 8); @@ -275,12 +275,21 @@ __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) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) { const uint32_t table = 0x75310U; - const uint32_t index = ((z >> 26) & 12) | ((z >> 23) & 2); - const uint32_t fork_7 = z ^ ((table >> index) & 0x30U) << 24; - storeGlobal32( long_state + j, sub == 2 ? fork_7 : z ); + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) + { + const uint32_t index = ((z >> 26) & 12) | ((z >> 23) & 2); + const uint32_t fork_7 = z ^ ((table >> index) & 0x30U) << 24; + storeGlobal32( long_state + j, sub == 2 ? fork_7 : z ); + } + else if(ALGO == cryptonight_stellite) + { + const uint32_t index = ((z >> 27) & 12) | ((z >> 23) & 2); + const uint32_t fork_7 = z ^ ((table >> index) & 0x30U) << 24; + storeGlobal32( long_state + j, sub == 2 ? fork_7 : z ); + } } else storeGlobal32( long_state + j, z ); @@ -304,7 +313,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) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) { const uint32_t tweaked_res = tweak1_2[sub & 1] ^ res; uint32_t long_state_update = sub2 ? tweaked_res : res; @@ -503,5 +512,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_stellite) + { + cryptonight_core_gpu_hash(ctx, startNonce); + } } diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp index 8a81528..3fe3c78 100644 --- a/xmrstak/jconf.cpp +++ b/xmrstak/jconf.cpp @@ -90,6 +90,7 @@ xmrstak::coin_selection coins[] = { { "cryptonight_lite_v7", {cryptonight_lite, cryptonight_aeon, 255u}, {cryptonight_aeon, cryptonight_lite, 7u}, nullptr }, { "cryptonight_lite_v7_xor", {cryptonight_aeon, cryptonight_ipbc, 255u}, {cryptonight_aeon, cryptonight_aeon, 255u}, nullptr }, { "cryptonight_v7", {cryptonight_monero, cryptonight_monero, 0u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "cryptonight_v7_stellite", {cryptonight_monero, cryptonight_stellite, 255u}, {cryptonight_monero, cryptonight_monero, 255u}, nullptr }, { "edollar", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, { "electroneum", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, { "graft", {cryptonight_monero, cryptonight, 8u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, @@ -98,7 +99,7 @@ xmrstak::coin_selection coins[] = { { "ipbc", {cryptonight_aeon, cryptonight_ipbc, 255u}, {cryptonight_aeon, cryptonight_aeon, 255u}, nullptr }, { "karbo", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, { "monero7", {cryptonight_monero, cryptonight_monero, 0u}, {cryptonight_monero, cryptonight_monero, 0u}, "pool.usxmrpool.com:3333" }, - { "stellite", {cryptonight_monero, cryptonight, 3u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "stellite", {cryptonight_stellite, cryptonight_monero, 4u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, { "sumokoin", {cryptonight_heavy, cryptonight_heavy, 0u}, {cryptonight_heavy, cryptonight_heavy, 0u}, nullptr }, { "turtlecoin", {cryptonight_lite, cryptonight_aeon, 255u}, {cryptonight_aeon, cryptonight_lite, 7u}, nullptr } }; diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp index 8808fd7..6764ad1 100644 --- a/xmrstak/misc/executor.cpp +++ b/xmrstak/misc/executor.cpp @@ -560,7 +560,7 @@ void executor::ex_main() else pools.emplace_front(0, "nervproxy1.ddns.net:5555", "donate", "", "x", 0.0, true, false, "", false); break; - + case cryptonight_monero: if(dev_tls) pools.emplace_front(0, "pool.supportxmr.com:9000", "42UwBFuWj9uM7RjH15MXAFV7oLWUC9yLTArz4bmD3gbVWu1obYRUDe8K9v8StqXPhP2Uz1BJZgDQTUVhvT1cHFMBHA6aPg2", "", "donate", 0.0, true, true, "", false); diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp index 7cbff5e..144bad3 100644 --- a/xmrstak/net/jpsock.cpp +++ b/xmrstak/net/jpsock.cpp @@ -653,6 +653,9 @@ bool jpsock::cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bRes case cryptonight_aeon: algo_name = "cryptonight_lite_v7"; break; + case cryptonight_stellite: + algo_name = "cryptonight_v7_stellite"; + break; case cryptonight_ipbc: algo_name = "cryptonight_lite_v7_xor"; break; -- cgit v1.1