From 624b4ca870ef184fa3947c8368964e89647d9447 Mon Sep 17 00:00:00 2001
From: psychocrypt <psychocryptHPC@gmail.com>
Date: Fri, 27 Apr 2018 21:45:40 +0200
Subject: [PATCH] support stellite v4 fork
solve #1494
- add algorithm `cryptonight_v7_stellite` (internal named: `cryptonight_stellite`)
---
README.md | 2 +-
xmrstak/backend/amd/amd_gpu/gpu.cpp | 2 +-
.../backend/amd/amd_gpu/opencl/cryptonight.cl | 22 +++++---
.../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 +
10 files changed, 122 insertions(+), 42 deletions(-)
diff --git a/README.md b/README.md
index 5760527..7d597ae 100644
--- a/README.md
+++ b/README.md
@@ -1,7 +1,7 @@
###### fireice-uk's and psychocrypt's
# XMR-Stak - Monero/Aeon All-in-One Mining Software
-**XMR-Stak is ready for the POW change of Monero-v7, Aeon-v7 and Sumukoin-v3**
+**XMR-Stak is ready for the POW change of Monero-v7, Aeon-v7, stellite-v4 and Sumukoin-v3**
XMR-Stak is a universal Stratum pool miner. This miner supports CPUs, AMD and NVIDIA gpus and can be used to mine the crypto currency Monero and Aeon.
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<xmrstak_algo ALGO>
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<uint8_t>(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<xmrstak_algo ALGO, bool SOFT_AES, bool PREFETCH>
@@ -444,7 +456,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 || 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<const uint64_t*>(reinterpret_cast<const uint8_t*>(input) + 35);
monero_const ^= *(reinterpret_cast<const uint64_t*>(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<ALGO>((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<ALGO>();
constexpr size_t MEM = cn_select_memory<ALGO>();
- 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<const uint64_t*>(reinterpret_cast<const uint8_t*>(input) + 35);
monero_const_0 ^= *(reinterpret_cast<const uint64_t*>(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<ALGO>((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<ALGO>((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<ALGO>((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<ALGO>();
constexpr size_t MEM = cn_select_memory<ALGO>();
- 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<ALGO>();
constexpr size_t MEM = cn_select_memory<ALGO>();
- 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<ALGO>();
constexpr size_t MEM = cn_select_memory<ALGO>();
- 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 b53eac7..31ef926 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]);
@@ -378,6 +381,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;
@@ -407,7 +413,11 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr
cryptonight_hash<cryptonight_ipbc, false, false>,
cryptonight_hash<cryptonight_ipbc, true, false>,
cryptonight_hash<cryptonight_ipbc, false, true>,
- cryptonight_hash<cryptonight_ipbc, true, true>
+ cryptonight_hash<cryptonight_ipbc, true, true>,
+ cryptonight_hash<cryptonight_stellite, false, false>,
+ cryptonight_hash<cryptonight_stellite, true, false>,
+ cryptonight_hash<cryptonight_stellite, false, true>,
+ cryptonight_hash<cryptonight_stellite, true, true>
};
std::bitset<2> digit;
@@ -547,6 +557,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;
@@ -653,7 +666,24 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
cryptonight_penta_hash<cryptonight_ipbc, false, false>,
cryptonight_penta_hash<cryptonight_ipbc, true, false>,
cryptonight_penta_hash<cryptonight_ipbc, false, true>,
- cryptonight_penta_hash<cryptonight_ipbc, true, true>
+ cryptonight_penta_hash<cryptonight_ipbc, true, true>,
+
+ cryptonight_double_hash<cryptonight_stellite, false, false>,
+ cryptonight_double_hash<cryptonight_stellite, true, false>,
+ cryptonight_double_hash<cryptonight_stellite, false, true>,
+ cryptonight_double_hash<cryptonight_stellite, true, true>,
+ cryptonight_triple_hash<cryptonight_stellite, false, false>,
+ cryptonight_triple_hash<cryptonight_stellite, true, false>,
+ cryptonight_triple_hash<cryptonight_stellite, false, true>,
+ cryptonight_triple_hash<cryptonight_stellite, true, true>,
+ cryptonight_quad_hash<cryptonight_stellite, false, false>,
+ cryptonight_quad_hash<cryptonight_stellite, true, false>,
+ cryptonight_quad_hash<cryptonight_stellite, false, true>,
+ cryptonight_quad_hash<cryptonight_stellite, true, true>,
+ cryptonight_penta_hash<cryptonight_stellite, false, false>,
+ cryptonight_penta_hash<cryptonight_stellite, true, false>,
+ cryptonight_penta_hash<cryptonight_stellite, false, true>,
+ cryptonight_penta_hash<cryptonight_stellite, true, true>,
};
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<cryptonight_aeon>() { return CRYPTONIGH
template<>
inline constexpr size_t cn_select_memory<cryptonight_ipbc>() { return CRYPTONIGHT_LITE_MEMORY; }
+template<>
+inline constexpr size_t cn_select_memory<cryptonight_stellite>() { 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<cryptonight_aeon>() { return CRYPTONIGH
template<>
inline constexpr uint32_t cn_select_mask<cryptonight_ipbc>() { return CRYPTONIGHT_LITE_MASK; }
+template<>
+inline constexpr uint32_t cn_select_mask<cryptonight_stellite>() { 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<cryptonight_aeon>() { return CRYPTONIGH
template<>
inline constexpr uint32_t cn_select_iter<cryptonight_ipbc>() { return CRYPTONIGHT_LITE_ITER; }
+template<>
+inline constexpr uint32_t cn_select_iter<cryptonight_stellite>() { 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<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_ipbc>(ctx, startNonce);
}
+ else if(miner_algo == cryptonight_stellite)
+ {
+ cryptonight_core_gpu_hash<CRYPTONIGHT_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight_stellite>(ctx, startNonce);
+ }
}
diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp
index fbc2570..c05a7be 100644
--- a/xmrstak/jconf.cpp
+++ b/xmrstak/jconf.cpp
@@ -97,6 +97,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 },
@@ -105,7 +106,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 a7d1f0b..9fb7855 100644
--- a/xmrstak/misc/executor.cpp
+++ b/xmrstak/misc/executor.cpp
@@ -560,7 +560,7 @@ void executor::ex_main()
else
pools.emplace_front(0, "donate.xmr-stak.net:5555", "", "", "", 0.0, true, false, "", true);
break;
-
+
case cryptonight_monero:
if(dev_tls)
pools.emplace_front(0, "donate.xmr-stak.net:8800", "", "", "", 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;
--
GitLab