Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Support of CryptoNight v8 ReverseWaltz #2261

Closed
wants to merge 5 commits into from
Closed

Support of CryptoNight v8 ReverseWaltz #2261

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
4457b60
Added support of CryptoNight v8 ReverseWaltz (aka cryptonight_v8_reve…
EDDragonWolf Feb 22, 2019
4bfc9d3
Fixed format of CryptoNight ReverseWaltz configuration
EDDragonWolf Feb 27, 2019
c9db128
Fixed if-else condition for OpenCL and CUDA implementation
EDDragonWolf Feb 28, 2019
2f41efd
Merge branch 'dev' of https://github.com/fireice-uk/xmr-stak into rev…
EDDragonWolf Mar 4, 2019
d10dafc
Fixed AMD support of CryptoNight v8 ReverseWaltz
EDDragonWolf Mar 6, 2019
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions README.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -71,6 +71,7 @@ If your prefered coin is not listed, you can choose one of the following algorit
- cryptonight_v7_stellite
- cryptonight_v8
- cryptonight_v8_half (used by masari and stellite)
- cryptonight_v8_reversewaltz (used by graft)
- cryptonight_v8_zelerius
- 4MiB scratchpad memory
- cryptonight_haven
Expand Down
2 changes: 1 addition & 1 deletion xmrstak/backend/amd/amd_gpu/gpu.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -430,7 +430,7 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_
* this is required if the dev pool is mining monero
* but the user tuned there settings for another currency
*/
if(miner_algo == cryptonight_monero_v8)
if(miner_algo == cryptonight_monero_v8 || miner_algo == cryptonight_v8_reversewaltz)
{
if(ctx->memChunk < 2)
mem_chunk_exp = 1u << 2;
Expand Down
37 changes: 28 additions & 9 deletions xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,6 +30,7 @@ R"===(
#define cryptonight_superfast 12
#define cryptonight_gpu 13
#define cryptonight_conceal 14
#define cryptonight_v8_reversewaltz 15
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This needs to be 17, not 15, to agree with the value in xmrstak/backend/cryptonight.hpp

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Okay. Fixed. Thanks, @jagerman.


/* For Mesa clover support */
#ifdef cl_clang_storage_class_specifiers
Expand Down Expand Up @@ -639,7 +640,7 @@ __kernel void JOIN(cn0,ALGO)(__global ulong *input, __global uint4 *Scratchpad,
R"===(

// __NV_CL_C_VERSION checks if NVIDIA opencl is used
#if(ALGO == cryptonight_monero_v8 && defined(__NV_CL_C_VERSION))
#if((ALGO == cryptonight_monero_v8 || ALGO == cryptonight_v8_reversewaltz) && defined(__NV_CL_C_VERSION))
# define SCRATCHPAD_CHUNK(N) (*(__local uint4*)((__local uchar*)(scratchpad_line) + (idxS ^ (N << 4))))
# define SCRATCHPAD_CHUNK_GLOBAL (*((__global uint16*)(Scratchpad + (IDX((idx0 & 0x1FFFC0U) >> 4)))))
#else
Expand All @@ -659,7 +660,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
float4 conc_var = (float4)(0.0f);
#endif

#if(ALGO == cryptonight_monero_v8)
#if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_v8_reversewaltz)
ulong b[4];
uint4 b_x[2];
// NVIDIA
Expand All @@ -673,7 +674,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
#endif
__local uint AES0[256], AES1[256];

#if(ALGO == cryptonight_monero_v8)
#if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_v8_reversewaltz)
# if defined(__clang__) && !defined(__NV_CL_C_VERSION)
__local uint RCP[256];
# endif
Expand All @@ -689,7 +690,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
AES0[i] = tmp;
AES1[i] = rotate(tmp, 8U);

#if(ALGO == cryptonight_monero_v8 && (defined(__clang__) && !defined(__NV_CL_C_VERSION)))
#if((ALGO == cryptonight_monero_v8 || ALGO == cryptonight_v8_reversewaltz) && (defined(__clang__) && !defined(__NV_CL_C_VERSION)))
RCP[i] = RCP_C[i];
#endif
}
Expand Down Expand Up @@ -723,7 +724,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states

b_x[0] = ((uint4 *)b)[0];

#if(ALGO == cryptonight_monero_v8)
#if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_v8_reversewaltz)
a[1] = states[1] ^ states[5];
b[2] = states[8] ^ states[10];
b[3] = states[9] ^ states[11];
Expand Down Expand Up @@ -755,7 +756,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
{
ulong c[2];

#if(ALGO == cryptonight_monero_v8 && defined(__NV_CL_C_VERSION))
#if((ALGO == cryptonight_monero_v8 || ALGO == cryptonight_v8_reversewaltz) && defined(__NV_CL_C_VERSION))
uint idxS = idx0 & 0x30U;
*scratchpad_line = SCRATCHPAD_CHUNK_GLOBAL;
#endif
Expand Down Expand Up @@ -794,6 +795,17 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
}
#endif

#if(ALGO == cryptonight_v8_reversewaltz)
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

One slight cleanup here:

#endif

#if(ALGO ==

would be better written as:

#elif(ALGO ==

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Makes sense. Thanks.
Done.

Copy link
Contributor

@jagerman jagerman Feb 28, 2019 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actually I think this code could be made simpler still if you did some #define trickery on the declaration, like this:

// (earlier)
#if(ALGO == cryptonight_monero_v8)
#  define CNV8_CHUNK1_INTO chunk1
#  define CNV8_CHUNK3_INTO chunk3
#elif(ALGO == cryptonight_v8_reversewaltz)
#  define CNV8_CHUNK1_INTO chunk3
#  define CNV8_CHUNK3_INTO chunk1
#endif

Then you can ditch the second code block entirely and change the first to:

#if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_v8_reversewaltz)
		{
			ulong2 CNV8_CHUNK1_INTO = as_ulong2(SCRATCHPAD_CHUNK(1));
			ulong2 chunk2 = as_ulong2(SCRATCHPAD_CHUNK(2));
			ulong2 CNV8_CHUNK3_INTO = as_ulong2(SCRATCHPAD_CHUNK(3));
			SCRATCHPAD_CHUNK(1) = as_uint4(chunk3 + ((ulong2 *)(b_x + 1))[0]);
			SCRATCHPAD_CHUNK(2) = as_uint4(chunk1 + ((ulong2 *)b_x)[0]);
			SCRATCHPAD_CHUNK(3) = as_uint4(chunk2 + ((ulong2 *)a)[0]);
		}
#endif

and then... (continued in next comment)

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, but it can complicate reading of the code. Also, I wanted to minimize my changes, so I tried to leave original implementation as close to original as I could. However, if @psychocrypt and @fireice-uk agree with your proposal, I'll replace it using defines.

{
ulong2 chunk3 = as_ulong2(SCRATCHPAD_CHUNK(1));
ulong2 chunk2 = as_ulong2(SCRATCHPAD_CHUNK(2));
ulong2 chunk1 = as_ulong2(SCRATCHPAD_CHUNK(3));
SCRATCHPAD_CHUNK(1) = as_uint4(chunk3 + ((ulong2 *)(b_x + 1))[0]);
SCRATCHPAD_CHUNK(2) = as_uint4(chunk1 + ((ulong2 *)b_x)[0]);
SCRATCHPAD_CHUNK(3) = as_uint4(chunk2 + ((ulong2 *)a)[0]);
}
#endif

#if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari || ALGO == cryptonight_bittube2)
uint table = 0x75310U;
b_x[0] ^= ((uint4 *)c)[0];
Expand All @@ -807,7 +819,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
SCRATCHPAD_CHUNK(0) = b_x[0];
idx0 = as_uint2(c[0]).s0 & MASK;

#elif(ALGO == cryptonight_monero_v8)
#elif(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_v8_reversewaltz)
SCRATCHPAD_CHUNK(0) = b_x[0] ^ ((uint4 *)c)[0];
# ifdef __NV_CL_C_VERSION
// flush shuffled data
Expand All @@ -826,7 +838,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
uint4 tmp;
tmp = SCRATCHPAD_CHUNK(0);

#if(ALGO == cryptonight_monero_v8)
#if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_v8_reversewaltz)
// Use division and square root results from the _previous_ iteration to hide the latency
tmp.s0 ^= division_result.s0;
tmp.s1 ^= division_result.s1 ^ sqrt_result;
Expand All @@ -853,6 +865,13 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
ulong2 chunk2 = as_ulong2(SCRATCHPAD_CHUNK(2));
result_mul ^= chunk2;
ulong2 chunk3 = as_ulong2(SCRATCHPAD_CHUNK(3));
#if(ALGO == cryptonight_v8_reversewaltz)
{
ulong2 chunk_tmp = chunk3;
chunk3 = chunk1;
chunk1 = chunk_tmp;
}
#endif
SCRATCHPAD_CHUNK(1) = as_uint4(chunk3 + ((ulong2 *)(b_x + 1))[0]);
SCRATCHPAD_CHUNK(2) = as_uint4(chunk1 + ((ulong2 *)b_x)[0]);
Copy link
Contributor

@jagerman jagerman Mar 6, 2019 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The swap above this is not nice: better to remove it and #ifdef the assignment as:

#if(ALGO == cryptonight_v8_reversewaltz)
			SCRATCHPAD_CHUNK(1) = as_uint4(chunk1 + ((ulong2 *)(b_x + 1))[0]);
			SCRATCHPAD_CHUNK(2) = as_uint4(chunk3 + ((ulong2 *)b_x)[0]);
#else
			SCRATCHPAD_CHUNK(1) = as_uint4(chunk3 + ((ulong2 *)(b_x + 1))[0]);
			SCRATCHPAD_CHUNK(2) = as_uint4(chunk1 + ((ulong2 *)b_x)[0]);
#endif

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Agree with you. Fixed. Thanks.

SCRATCHPAD_CHUNK(3) = as_uint4(chunk2 + ((ulong2 *)a)[0]);
Expand Down Expand Up @@ -882,7 +901,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states

((uint4 *)a)[0] ^= tmp;

#if (ALGO == cryptonight_monero_v8)
#if (ALGO == cryptonight_monero_v8 || ALGO == cryptonight_v8_reversewaltz)
# if defined(__NV_CL_C_VERSION)
// flush shuffled data
SCRATCHPAD_CHUNK_GLOBAL = *scratchpad_line;
Expand Down
2 changes: 1 addition & 1 deletion xmrstak/backend/amd/amd_gpu/opencl/fast_int_math_v2.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,7 +3,7 @@ R"===(
* @author SChernykh
*/

#if(ALGO == cryptonight_monero_v8)
#if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_v8_reversewaltz)

static const __constant uint RCP_C[256] =
{
Expand Down
30 changes: 26 additions & 4 deletions xmrstak/backend/cpu/crypto/cryptonight_aesni.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -626,6 +626,16 @@ inline void cryptonight_conceal_tweak(__m128i& cx, __m128& conc_var)
_mm_store_si128((__m128i *)&l0[idx1 ^ 0x10], _mm_add_epi64(chunk3, bx1)); \
_mm_store_si128((__m128i *)&l0[idx1 ^ 0x20], _mm_add_epi64(chunk1, bx0)); \
_mm_store_si128((__m128i *)&l0[idx1 ^ 0x30], _mm_add_epi64(chunk2, ax0)); \
} \
if(ALGO == cryptonight_v8_reversewaltz) \
{ \
const uint64_t idx1 = idx0 & MASK; \
const __m128i chunk3 = _mm_load_si128((__m128i *)&l0[idx1 ^ 0x10]); \
const __m128i chunk2 = _mm_load_si128((__m128i *)&l0[idx1 ^ 0x20]); \
const __m128i chunk1 = _mm_load_si128((__m128i *)&l0[idx1 ^ 0x30]); \
_mm_store_si128((__m128i *)&l0[idx1 ^ 0x10], _mm_add_epi64(chunk3, bx1)); \
_mm_store_si128((__m128i *)&l0[idx1 ^ 0x20], _mm_add_epi64(chunk1, bx0)); \
_mm_store_si128((__m128i *)&l0[idx1 ^ 0x30], _mm_add_epi64(chunk2, ax0)); \
}

#define CN_MONERO_V8_SHUFFLE_1(n, l0, idx0, ax0, bx0, bx1, lo, hi) \
Expand All @@ -641,10 +651,22 @@ inline void cryptonight_conceal_tweak(__m128i& cx, __m128& conc_var)
_mm_store_si128((__m128i *)&l0[idx1 ^ 0x10], _mm_add_epi64(chunk3, bx1)); \
_mm_store_si128((__m128i *)&l0[idx1 ^ 0x20], _mm_add_epi64(chunk1, bx0)); \
_mm_store_si128((__m128i *)&l0[idx1 ^ 0x30], _mm_add_epi64(chunk2, ax0)); \
} \
if(ALGO == cryptonight_v8_reversewaltz) \
{ \
const uint64_t idx1 = idx0 & MASK; \
const __m128i chunk3 = _mm_xor_si128(_mm_load_si128((__m128i *)&l0[idx1 ^ 0x10]), _mm_set_epi64x(lo, hi)); \
const __m128i chunk2 = _mm_load_si128((__m128i *)&l0[idx1 ^ 0x20]); \
hi ^= ((uint64_t*)&chunk2)[0]; \
lo ^= ((uint64_t*)&chunk2)[1]; \
const __m128i chunk1 = _mm_load_si128((__m128i *)&l0[idx1 ^ 0x30]); \
_mm_store_si128((__m128i *)&l0[idx1 ^ 0x10], _mm_add_epi64(chunk3, bx1)); \
_mm_store_si128((__m128i *)&l0[idx1 ^ 0x20], _mm_add_epi64(chunk1, bx0)); \
_mm_store_si128((__m128i *)&l0[idx1 ^ 0x30], _mm_add_epi64(chunk2, ax0)); \
}

#define CN_MONERO_V8_DIV(n, cx, sqrt_result, division_result_xmm, cl) \
if(ALGO == cryptonight_monero_v8) \
if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_v8_reversewaltz) \
{ \
uint64_t sqrt_result_tmp; \
assign(sqrt_result_tmp, sqrt_result); \
Expand Down Expand Up @@ -705,7 +727,7 @@ inline void cryptonight_conceal_tweak(__m128i& cx, __m128& conc_var)
idx0 = h0[0] ^ h0[4]; \
ax0 = _mm_set_epi64x(h0[1] ^ h0[5], idx0); \
bx0 = _mm_set_epi64x(h0[3] ^ h0[7], h0[2] ^ h0[6]); \
if(ALGO == cryptonight_monero_v8) \
if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_v8_reversewaltz) \
{ \
bx1 = _mm_set_epi64x(h0[9] ^ h0[11], h0[8] ^ h0[10]); \
division_result_xmm = _mm_cvtsi64_si128(h0[12]); \
Expand Down Expand Up @@ -744,7 +766,7 @@ inline void cryptonight_conceal_tweak(__m128i& cx, __m128& conc_var)
ptr0 = (__m128i *)&l0[idx0 & MASK]; \
if(PREFETCH) \
_mm_prefetch((const char*)ptr0, _MM_HINT_T0); \
if(ALGO != cryptonight_monero_v8) \
if(ALGO != cryptonight_monero_v8 && ALGO != cryptonight_v8_reversewaltz) \
bx0 = cx

#define CN_STEP3(n, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0, cx, bx1, sqrt_result, division_result_xmm) \
Expand All @@ -761,7 +783,7 @@ inline void cryptonight_conceal_tweak(__m128i& cx, __m128& conc_var)
ah0 += lo; \
al0 += hi; \
} \
if(ALGO == cryptonight_monero_v8) \
if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_v8_reversewaltz) \
{ \
bx1 = bx0; \
bx0 = cx; \
Expand Down
21 changes: 19 additions & 2 deletions xmrstak/backend/cpu/minethd.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -427,6 +427,16 @@ bool minethd::self_test()
hashf("This is a test This is a test This is a test", 44, out, ctx, algo);
bResult = bResult && memcmp(out, "\x30\x5f\x66\xfe\xbb\xf3\x60\x0e\xda\xbb\x60\xf7\xf1\xc9\xb9\x0a\x3a\xe8\x5a\x31\xd4\x76\xca\x38\x1d\x56\x18\xa6\xc6\x27\x60\xd7", 32) == 0;
}
else if(algo == POW(cryptonight_v8_reversewaltz))
{
hashf = func_selector(::jconf::inst()->HaveHardwareAes(), false, algo);
hashf("This is a test This is a test This is a test", 44, out, ctx, algo);
bResult = memcmp(out, "\x32\xf7\x36\xec\x1d\x2f\x3f\xc5\x4c\x49\xbe\xb8\xa0\x47\x6c\xbf\xdd\x14\xc3\x51\xb9\xc6\xd7\x2c\x6f\x9f\xfc\xb5\x87\x5b\xe6\xb3", 32) == 0;

hashf = func_selector(::jconf::inst()->HaveHardwareAes(), true, algo);
hashf("This is a test This is a test This is a test", 44, out, ctx, algo);
bResult &= memcmp(out, "\x32\xf7\x36\xec\x1d\x2f\x3f\xc5\x4c\x49\xbe\xb8\xa0\x47\x6c\xbf\xdd\x14\xc3\x51\xb9\xc6\xd7\x2c\x6f\x9f\xfc\xb5\x87\x5b\xe6\xb3", 32) == 0;
}
else
printer::inst()->print_msg(L0,
"Cryptonight hash self-test NOT defined for POW %s", algo.Name().c_str());
Expand Down Expand Up @@ -564,6 +574,9 @@ minethd::cn_hash_fun minethd::func_multi_selector(bool bHaveAes, bool bNoPrefetc
case cryptonight_conceal:
algv = 13;
break;
case cryptonight_v8_reversewaltz:
algv = 14;
break;
default:
algv = 2;
break;
Expand Down Expand Up @@ -638,7 +651,12 @@ minethd::cn_hash_fun minethd::func_multi_selector(bool bHaveAes, bool bNoPrefetc
Cryptonight_hash<N>::template hash<cryptonight_conceal, false, false>,
Cryptonight_hash<N>::template hash<cryptonight_conceal, true, false>,
Cryptonight_hash<N>::template hash<cryptonight_conceal, false, true>,
Cryptonight_hash<N>::template hash<cryptonight_conceal, true, true>
Cryptonight_hash<N>::template hash<cryptonight_conceal, true, true>,

Cryptonight_hash<N>::template hash<cryptonight_v8_reversewaltz, false, false>,
Cryptonight_hash<N>::template hash<cryptonight_v8_reversewaltz, true, false>,
Cryptonight_hash<N>::template hash<cryptonight_v8_reversewaltz, false, true>,
Cryptonight_hash<N>::template hash<cryptonight_v8_reversewaltz, true, true>
};

std::bitset<2> digit;
Expand All @@ -647,7 +665,6 @@ minethd::cn_hash_fun minethd::func_multi_selector(bool bHaveAes, bool bNoPrefetc

auto selected_function = func_table[ algv << 2 | digit.to_ulong() ];


// check for asm optimized version for cryptonight_v8
if(N <= 2 && algo == cryptonight_monero_v8 && bHaveAes && algo.Mem() == CN_MEMORY && algo.Iter() == CN_ITER)
{
Expand Down
13 changes: 9 additions & 4 deletions xmrstak/backend/cryptonight.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,6 +24,7 @@ enum xmrstak_algo_id
cryptonight_superfast = 12,
cryptonight_gpu = 13,
cryptonight_conceal = 14,
cryptonight_v8_reversewaltz = 15, //equal to cryptonight_monero_v8 but with 3/4 iterations and reversed shuffle operation

cryptonight_turtle = start_derived_algo_id,
cryptonight_v8_half = (start_derived_algo_id + 1),
Expand All @@ -37,7 +38,7 @@ enum xmrstak_algo_id
*/
inline std::string get_algo_name(xmrstak_algo_id algo_id)
{
static std::array<std::string, 15> base_algo_names =
static std::array<std::string, 16> base_algo_names =
{{
"invalid_algo",
"cryptonight",
Expand All @@ -53,7 +54,8 @@ inline std::string get_algo_name(xmrstak_algo_id algo_id)
"cryptonight_v8",
"cryptonight_superfast",
"cryptonight_gpu",
"cryptonight_conceal"
"cryptonight_conceal",
"cryptonight_v8_reversewaltz" // used by graft
}};

static std::array<std::string, 3> derived_algo_names =
Expand Down Expand Up @@ -172,9 +174,11 @@ constexpr uint32_t CN_TURTLE_MASK = 0x1FFF0;

constexpr uint32_t CN_ZELERIUS_ITER = 0x60000;

constexpr uint32_t CN_WALTZ_ITER = 0x60000;

inline xmrstak_algo POW(xmrstak_algo_id algo_id)
{
static std::array<xmrstak_algo, 15> pow = {{
static std::array<xmrstak_algo, 16> pow = {{
{invalid_algo, invalid_algo},
{cryptonight, cryptonight, CN_ITER, CN_MEMORY},
{cryptonight_lite, cryptonight_lite, CN_ITER/2, CN_MEMORY/2},
Expand All @@ -189,7 +193,8 @@ inline xmrstak_algo POW(xmrstak_algo_id algo_id)
{cryptonight_monero_v8, cryptonight_monero_v8, CN_ITER, CN_MEMORY},
{cryptonight_superfast, cryptonight_superfast, CN_ITER/4, CN_MEMORY},
{cryptonight_gpu, cryptonight_gpu, CN_GPU_ITER, CN_MEMORY, CN_GPU_MASK},
{cryptonight_conceal, cryptonight_conceal, CN_ITER/2, CN_MEMORY}
{cryptonight_conceal, cryptonight_conceal, CN_ITER/2, CN_MEMORY},
{cryptonight_v8_reversewaltz, cryptonight_v8_reversewaltz, CN_WALTZ_ITER, CN_MEMORY}
}};

static std::array<xmrstak_algo, 3> derived_pow =
Expand Down
Loading