diff --git a/src/base/net/Pool.cpp b/src/base/net/Pool.cpp index f66a8c9b..4a624c28 100644 --- a/src/base/net/Pool.cpp +++ b/src/base/net/Pool.cpp @@ -479,6 +479,8 @@ void xmrig::Pool::rebuild() m_algorithms.push_back(m_algorithm); # ifndef XMRIG_PROXY_PROJECT + addVariant(VARIANT_RWZ); + addVariant(VARIANT_WALTZ); addVariant(VARIANT_4); addVariant(VARIANT_WOW); addVariant(VARIANT_2); diff --git a/src/common/crypto/Algorithm.cpp b/src/common/crypto/Algorithm.cpp index dd864705..0a4e3953 100644 --- a/src/common/crypto/Algorithm.cpp +++ b/src/common/crypto/Algorithm.cpp @@ -54,18 +54,20 @@ struct AlgoData static AlgoData const algorithms[] = { - { "cryptonight", "cn", xmrig::CRYPTONIGHT, xmrig::VARIANT_AUTO }, - { "cryptonight/0", "cn/0", xmrig::CRYPTONIGHT, xmrig::VARIANT_0 }, - { "cryptonight/1", "cn/1", xmrig::CRYPTONIGHT, xmrig::VARIANT_1 }, - { "cryptonight/xtl", "cn/xtl", xmrig::CRYPTONIGHT, xmrig::VARIANT_XTL }, - { "cryptonight/msr", "cn/msr", xmrig::CRYPTONIGHT, xmrig::VARIANT_MSR }, - { "cryptonight/xao", "cn/xao", xmrig::CRYPTONIGHT, xmrig::VARIANT_XAO }, - { "cryptonight/rto", "cn/rto", xmrig::CRYPTONIGHT, xmrig::VARIANT_RTO }, - { "cryptonight/2", "cn/2", xmrig::CRYPTONIGHT, xmrig::VARIANT_2 }, - { "cryptonight/half", "cn/half", xmrig::CRYPTONIGHT, xmrig::VARIANT_HALF }, - { "cryptonight/xtlv9", "cn/xtlv9", xmrig::CRYPTONIGHT, xmrig::VARIANT_HALF }, - { "cryptonight/wow", "cn/wow", xmrig::CRYPTONIGHT, xmrig::VARIANT_WOW }, - { "cryptonight/r", "cn/r", xmrig::CRYPTONIGHT, xmrig::VARIANT_4 }, + { "cryptonight", "cn", xmrig::CRYPTONIGHT, xmrig::VARIANT_AUTO }, + { "cryptonight/0", "cn/0", xmrig::CRYPTONIGHT, xmrig::VARIANT_0 }, + { "cryptonight/1", "cn/1", xmrig::CRYPTONIGHT, xmrig::VARIANT_1 }, + { "cryptonight/xtl", "cn/xtl", xmrig::CRYPTONIGHT, xmrig::VARIANT_XTL }, + { "cryptonight/msr", "cn/msr", xmrig::CRYPTONIGHT, xmrig::VARIANT_MSR }, + { "cryptonight/xao", "cn/xao", xmrig::CRYPTONIGHT, xmrig::VARIANT_XAO }, + { "cryptonight/rto", "cn/rto", xmrig::CRYPTONIGHT, xmrig::VARIANT_RTO }, + { "cryptonight/2", "cn/2", xmrig::CRYPTONIGHT, xmrig::VARIANT_2 }, + { "cryptonight/half", "cn/half", xmrig::CRYPTONIGHT, xmrig::VARIANT_HALF }, + { "cryptonight/xtlv9", "cn/xtlv9", xmrig::CRYPTONIGHT, xmrig::VARIANT_HALF }, + { "cryptonight/wow", "cn/wow", xmrig::CRYPTONIGHT, xmrig::VARIANT_WOW }, + { "cryptonight/r", "cn/r", xmrig::CRYPTONIGHT, xmrig::VARIANT_4 }, + { "cryptonight/waltz", "cn/waltz", xmrig::CRYPTONIGHT, xmrig::VARIANT_WALTZ }, + { "cryptonight/rwz", "cn/rwz", xmrig::CRYPTONIGHT, xmrig::VARIANT_RWZ }, # ifndef XMRIG_NO_AEON { "cryptonight-lite", "cn-lite", xmrig::CRYPTONIGHT_LITE, xmrig::VARIANT_AUTO }, @@ -132,6 +134,8 @@ static const char *variants[] = { "gpu", "wow", "r", + "waltz", + "rwz" }; @@ -193,7 +197,6 @@ void xmrig::Algorithm::parseAlgorithm(const char *algo) } } - void xmrig::Algorithm::parseVariant(const char *variant) { m_variant = VARIANT_AUTO; diff --git a/src/common/xmrig.h b/src/common/xmrig.h index c6a5f568..e5cd4838 100644 --- a/src/common/xmrig.h +++ b/src/common/xmrig.h @@ -61,21 +61,23 @@ enum AlgoVariant { enum Variant { - VARIANT_AUTO = -1, // Autodetect - VARIANT_0 = 0, // Original CryptoNight or CryptoNight-Heavy - VARIANT_1 = 1, // CryptoNight variant 1 also known as Monero7 and CryptoNightV7 - VARIANT_TUBE = 2, // Modified CryptoNight-Heavy (TUBE only) - VARIANT_XTL = 3, // Modified CryptoNight variant 1 (Stellite only) - VARIANT_MSR = 4, // Modified CryptoNight variant 1 (Masari only) - VARIANT_XHV = 5, // Modified CryptoNight-Heavy (Haven Protocol only) - VARIANT_XAO = 6, // Modified CryptoNight variant 0 (Alloy only) - VARIANT_RTO = 7, // Modified CryptoNight variant 1 (Arto only) - VARIANT_2 = 8, // CryptoNight variant 2 - VARIANT_HALF = 9, // CryptoNight variant 2 with half iterations (Masari/Stellite) - VARIANT_TRTL = 10, // CryptoNight Turtle (TRTL) - VARIANT_GPU = 11, // CryptoNight-GPU (Ryo) - VARIANT_WOW = 12, // CryptoNightR (Wownero) - VARIANT_4 = 13, // CryptoNightR (Monero's variant 4) + VARIANT_AUTO = -1, // Autodetect + VARIANT_0 = 0, // Original CryptoNight or CryptoNight-Heavy + VARIANT_1 = 1, // CryptoNight variant 1 also known as Monero7 and CryptoNightV7 + VARIANT_TUBE = 2, // Modified CryptoNight-Heavy (TUBE only) + VARIANT_XTL = 3, // Modified CryptoNight variant 1 (Stellite only) + VARIANT_MSR = 4, // Modified CryptoNight variant 1 (Masari only) + VARIANT_XHV = 5, // Modified CryptoNight-Heavy (Haven Protocol only) + VARIANT_XAO = 6, // Modified CryptoNight variant 0 (Alloy only) + VARIANT_RTO = 7, // Modified CryptoNight variant 1 (Arto only) + VARIANT_2 = 8, // CryptoNight variant 2 + VARIANT_HALF = 9, // CryptoNight variant 2 with half iterations (Masari/Stellite) + VARIANT_TRTL = 10, // CryptoNight Turtle (TRTL) + VARIANT_GPU = 11, // CryptoNight-GPU (Ryo) + VARIANT_WOW = 12, // CryptoNightR (Wownero) + VARIANT_4 = 13, // CryptoNightR (Monero's variant 4) + VARIANT_WALTZ = 14, // CryptoNight variant 2 with 3/4 iterations + VARIANT_RWZ = 15, // CryptoNight variant 2 with 3/4 iterations and reversed shuffle operation (Graft) VARIANT_MAX }; diff --git a/src/crypto/CryptoNight_arm.h b/src/crypto/CryptoNight_arm.h index e7232eb1..d762929c 100644 --- a/src/crypto/CryptoNight_arm.h +++ b/src/crypto/CryptoNight_arm.h @@ -436,7 +436,7 @@ static inline void cryptonight_monero_tweak(const uint8_t* l, uint64_t idx, __m1 uint64_t* mem_out = (uint64_t*)&l[idx]; if (BASE == xmrig::VARIANT_2) { - VARIANT2_SHUFFLE(l, idx, ax0, bx0, bx1, cx); + VARIANT2_SHUFFLE(l, idx, ax0, bx0, bx1, cx, (VARIANT == xmrig::VARIANT_RWZ ? 1 : 0)); _mm_store_si128((__m128i *)mem_out, _mm_xor_si128(bx0, cx)); } else { __m128i tmp = _mm_xor_si128(bx0, cx); @@ -530,9 +530,9 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si if (BASE == xmrig::VARIANT_2) { if (VARIANT == xmrig::VARIANT_4) { - VARIANT2_SHUFFLE(l0, idx0 & MASK, ax0, bx0, bx1, cx); + VARIANT2_SHUFFLE(l0, idx0 & MASK, ax0, bx0, bx1, cx, 0); } else { - VARIANT2_SHUFFLE2(l0, idx0 & MASK, ax0, bx0, bx1, hi, lo); + VARIANT2_SHUFFLE2(l0, idx0 & MASK, ax0, bx0, bx1, hi, lo, (VARIANT == xmrig::VARIANT_RWZ ? 1 : 0)); } } @@ -709,9 +709,9 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si if (BASE == xmrig::VARIANT_2) { if (VARIANT == xmrig::VARIANT_4) { - VARIANT2_SHUFFLE(l0, idx0 & MASK, ax0, bx00, bx01, cx0); + VARIANT2_SHUFFLE(l0, idx0 & MASK, ax0, bx00, bx01, cx0, 0); } else { - VARIANT2_SHUFFLE2(l0, idx0 & MASK, ax0, bx00, bx01, hi, lo); + VARIANT2_SHUFFLE2(l0, idx0 & MASK, ax0, bx00, bx01, hi, lo, (VARIANT == xmrig::VARIANT_RWZ ? 1 : 0)); } } @@ -767,9 +767,9 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si if (BASE == xmrig::VARIANT_2) { if (VARIANT == xmrig::VARIANT_4) { - VARIANT2_SHUFFLE(l1, idx1 & MASK, ax1, bx10, bx11, cx1); + VARIANT2_SHUFFLE(l1, idx1 & MASK, ax1, bx10, bx11, cx1, 0); } else { - VARIANT2_SHUFFLE2(l1, idx1 & MASK, ax1, bx10, bx11, hi, lo); + VARIANT2_SHUFFLE2(l1, idx1 & MASK, ax1, bx10, bx11, hi, lo, (VARIANT == xmrig::VARIANT_RWZ ? 1 : 0)); } } diff --git a/src/crypto/CryptoNight_constants.h b/src/crypto/CryptoNight_constants.h index 4ea1adb3..581c0c5a 100644 --- a/src/crypto/CryptoNight_constants.h +++ b/src/crypto/CryptoNight_constants.h @@ -42,6 +42,7 @@ constexpr const uint32_t CRYPTONIGHT_MASK = 0x1FFFF0; constexpr const uint32_t CRYPTONIGHT_ITER = 0x80000; constexpr const uint32_t CRYPTONIGHT_HALF_ITER = 0x40000; constexpr const uint32_t CRYPTONIGHT_XAO_ITER = 0x100000; +constexpr const uint32_t CRYPTONIGHT_WALTZ_ITER = 0x60000; constexpr const uint32_t CRYPTONIGHT_GPU_ITER = 0xC000; constexpr const uint32_t CRYPTONIGHT_GPU_MASK = 0x1FFFC0; @@ -134,6 +135,8 @@ template<> inline constexpr uint32_t cn_select_iter() template<> inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_XAO_ITER; } template<> inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_ITER; } template<> inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_GPU_ITER; } +template<> inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_WALTZ_ITER; } +template<> inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_WALTZ_ITER; } template<> inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_LITE_ITER; } template<> inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_LITE_ITER; } template<> inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_HEAVY_ITER; } @@ -158,6 +161,10 @@ inline uint32_t cn_select_iter(Algo algorithm, Variant variant) case VARIANT_TRTL: return CRYPTONIGHT_TRTL_ITER; + case VARIANT_WALTZ: + case VARIANT_RWZ: + return CRYPTONIGHT_WALTZ_ITER; + default: break; } @@ -199,6 +206,8 @@ template<> inline constexpr Variant cn_base_variant() { return VA template<> inline constexpr Variant cn_base_variant() { return VARIANT_GPU; } template<> inline constexpr Variant cn_base_variant() { return VARIANT_2; } template<> inline constexpr Variant cn_base_variant() { return VARIANT_2; } +template<> inline constexpr Variant cn_base_variant() { return VARIANT_2; } +template<> inline constexpr Variant cn_base_variant() { return VARIANT_2; } template inline constexpr bool cn_is_cryptonight_r() { return false; } diff --git a/src/crypto/CryptoNight_monero.h b/src/crypto/CryptoNight_monero.h index 26c1fff0..f2310b19 100644 --- a/src/crypto/CryptoNight_monero.h +++ b/src/crypto/CryptoNight_monero.h @@ -83,11 +83,11 @@ sqrt_result_xmm_##part = int_sqrt_v2(cx_0 + division_result); \ } while (0) -# define VARIANT2_SHUFFLE(base_ptr, offset, _a, _b, _b1, _c) \ +# define VARIANT2_SHUFFLE(base_ptr, offset, _a, _b, _b1, _c, reverse) \ do { \ - const __m128i chunk1 = _mm_load_si128((__m128i *)((base_ptr) + ((offset) ^ 0x10))); \ + const __m128i chunk1 = _mm_load_si128((__m128i *)((base_ptr) + ((offset) ^ (reverse ? 0x30 : 0x10)))); \ const __m128i chunk2 = _mm_load_si128((__m128i *)((base_ptr) + ((offset) ^ 0x20))); \ - const __m128i chunk3 = _mm_load_si128((__m128i *)((base_ptr) + ((offset) ^ 0x30))); \ + const __m128i chunk3 = _mm_load_si128((__m128i *)((base_ptr) + ((offset) ^ (reverse ? 0x10 : 0x30)))); \ _mm_store_si128((__m128i *)((base_ptr) + ((offset) ^ 0x10)), _mm_add_epi64(chunk3, _b1)); \ _mm_store_si128((__m128i *)((base_ptr) + ((offset) ^ 0x20)), _mm_add_epi64(chunk1, _b)); \ _mm_store_si128((__m128i *)((base_ptr) + ((offset) ^ 0x30)), _mm_add_epi64(chunk2, _a)); \ @@ -96,15 +96,20 @@ } \ } while (0) -# define VARIANT2_SHUFFLE2(base_ptr, offset, _a, _b, _b1, hi, lo) \ +# define VARIANT2_SHUFFLE2(base_ptr, offset, _a, _b, _b1, hi, lo, reverse) \ do { \ const __m128i chunk1 = _mm_xor_si128(_mm_load_si128((__m128i *)((base_ptr) + ((offset) ^ 0x10))), _mm_set_epi64x(lo, hi)); \ const __m128i chunk2 = _mm_load_si128((__m128i *)((base_ptr) + ((offset) ^ 0x20))); \ hi ^= ((uint64_t*)((base_ptr) + ((offset) ^ 0x20)))[0]; \ lo ^= ((uint64_t*)((base_ptr) + ((offset) ^ 0x20)))[1]; \ const __m128i chunk3 = _mm_load_si128((__m128i *)((base_ptr) + ((offset) ^ 0x30))); \ - _mm_store_si128((__m128i *)((base_ptr) + ((offset) ^ 0x10)), _mm_add_epi64(chunk3, _b1)); \ - _mm_store_si128((__m128i *)((base_ptr) + ((offset) ^ 0x20)), _mm_add_epi64(chunk1, _b)); \ + if (reverse) { \ + _mm_store_si128((__m128i *)((base_ptr) + ((offset) ^ 0x10)), _mm_add_epi64(chunk1, _b1)); \ + _mm_store_si128((__m128i *)((base_ptr) + ((offset) ^ 0x20)), _mm_add_epi64(chunk3, _b)); \ + } else { \ + _mm_store_si128((__m128i *)((base_ptr) + ((offset) ^ 0x10)), _mm_add_epi64(chunk3, _b1)); \ + _mm_store_si128((__m128i *)((base_ptr) + ((offset) ^ 0x20)), _mm_add_epi64(chunk1, _b)); \ + } \ _mm_store_si128((__m128i *)((base_ptr) + ((offset) ^ 0x30)), _mm_add_epi64(chunk2, _a)); \ } while (0) @@ -128,11 +133,11 @@ sqrt_result_##part += ((r2 + b > sqrt_input) ? -1 : 0) + ((r2 + (1ULL << 32) < sqrt_input - s) ? 1 : 0); \ } while (0) -# define VARIANT2_SHUFFLE(base_ptr, offset, _a, _b, _b1, _c) \ +# define VARIANT2_SHUFFLE(base_ptr, offset, _a, _b, _b1, _c, reverse) \ do { \ - const uint64x2_t chunk1 = vld1q_u64((uint64_t*)((base_ptr) + ((offset) ^ 0x10))); \ + const uint64x2_t chunk1 = vld1q_u64((uint64_t*)((base_ptr) + ((offset) ^ (reverse ? 0x30 : 0x10)))); \ const uint64x2_t chunk2 = vld1q_u64((uint64_t*)((base_ptr) + ((offset) ^ 0x20))); \ - const uint64x2_t chunk3 = vld1q_u64((uint64_t*)((base_ptr) + ((offset) ^ 0x30))); \ + const uint64x2_t chunk3 = vld1q_u64((uint64_t*)((base_ptr) + ((offset) ^ (reverse ? 0x10 : 0x30)))); \ vst1q_u64((uint64_t*)((base_ptr) + ((offset) ^ 0x10)), vaddq_u64(chunk3, vreinterpretq_u64_u8(_b1))); \ vst1q_u64((uint64_t*)((base_ptr) + ((offset) ^ 0x20)), vaddq_u64(chunk1, vreinterpretq_u64_u8(_b))); \ vst1q_u64((uint64_t*)((base_ptr) + ((offset) ^ 0x30)), vaddq_u64(chunk2, vreinterpretq_u64_u8(_a))); \ @@ -141,15 +146,20 @@ } \ } while (0) -# define VARIANT2_SHUFFLE2(base_ptr, offset, _a, _b, _b1, hi, lo) \ +# define VARIANT2_SHUFFLE2(base_ptr, offset, _a, _b, _b1, hi, lo, reverse) \ do { \ const uint64x2_t chunk1 = veorq_u64(vld1q_u64((uint64_t*)((base_ptr) + ((offset) ^ 0x10))), vcombine_u64(vcreate_u64(hi), vcreate_u64(lo))); \ const uint64x2_t chunk2 = vld1q_u64((uint64_t*)((base_ptr) + ((offset) ^ 0x20))); \ hi ^= ((uint64_t*)((base_ptr) + ((offset) ^ 0x20)))[0]; \ lo ^= ((uint64_t*)((base_ptr) + ((offset) ^ 0x20)))[1]; \ const uint64x2_t chunk3 = vld1q_u64((uint64_t*)((base_ptr) + ((offset) ^ 0x30))); \ - vst1q_u64((uint64_t*)((base_ptr) + ((offset) ^ 0x10)), vaddq_u64(chunk3, vreinterpretq_u64_u8(_b1))); \ - vst1q_u64((uint64_t*)((base_ptr) + ((offset) ^ 0x20)), vaddq_u64(chunk1, vreinterpretq_u64_u8(_b))); \ + if (reverse) { \ + vst1q_u64((uint64_t*)((base_ptr) + ((offset) ^ 0x10)), vaddq_u64(chunk1, vreinterpretq_u64_u8(_b1))); \ + vst1q_u64((uint64_t*)((base_ptr) + ((offset) ^ 0x20)), vaddq_u64(chunk3, vreinterpretq_u64_u8(_b))); \ + } else { \ + vst1q_u64((uint64_t*)((base_ptr) + ((offset) ^ 0x10)), vaddq_u64(chunk3, vreinterpretq_u64_u8(_b1))); \ + vst1q_u64((uint64_t*)((base_ptr) + ((offset) ^ 0x20)), vaddq_u64(chunk1, vreinterpretq_u64_u8(_b))); \ + } \ vst1q_u64((uint64_t*)((base_ptr) + ((offset) ^ 0x30)), vaddq_u64(chunk2, vreinterpretq_u64_u8(_a))); \ } while (0) #endif diff --git a/src/crypto/CryptoNight_test.h b/src/crypto/CryptoNight_test.h index 237fe31b..6fb0082b 100644 --- a/src/crypto/CryptoNight_test.h +++ b/src/crypto/CryptoNight_test.h @@ -199,6 +199,32 @@ const static uint8_t test_output_rto[160] = { 0xE7, 0x81, 0x4E, 0x2A, 0xBD, 0x62, 0xC1, 0x1B, 0x7C, 0xB9, 0x33, 0x7B, 0xEE, 0x95, 0x80, 0xB3 }; +const static uint8_t test_output_waltz[160] = { + 0x51, 0x6e, 0x33, 0xc6, 0xe4, 0x46, 0xab, 0xbc, 0xcd, 0xad, 0x18, 0xc0, 0x4c, 0xd9, 0xa2, 0x5e, + 0x64, 0x10, 0x28, 0x53, 0xb2, 0x0a, 0x42, 0xdf, 0xde, 0xaa, 0x8b, 0x59, 0x9e, 0xcf, 0x40, 0xe2, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, +}; + +const static uint8_t test_output_rwz[160] = { + 0x5f, 0x56, 0xc6, 0xb0, 0x99, 0x6b, 0xa2, 0x3e, 0x0b, 0xba, 0x07, 0x29, 0xc9, 0x90, 0x74, 0x85, + 0x5a, 0x10, 0xe3, 0x08, 0x7f, 0xdb, 0xfe, 0x94, 0x75, 0x33, 0x54, 0x73, 0x76, 0xf0, 0x75, 0xb8, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, +}; + #ifndef XMRIG_NO_AEON // "cn-lite/0" diff --git a/src/crypto/CryptoNight_x86.h b/src/crypto/CryptoNight_x86.h index 4c5d4ac0..c6bfeed2 100644 --- a/src/crypto/CryptoNight_x86.h +++ b/src/crypto/CryptoNight_x86.h @@ -460,7 +460,7 @@ template static inline void cryptonight_monero_tweak(uint64_t* mem_out, const uint8_t* l, uint64_t idx, __m128i ax0, __m128i bx0, __m128i bx1, __m128i& cx) { if (BASE == xmrig::VARIANT_2) { - VARIANT2_SHUFFLE(l, idx, ax0, bx0, bx1, cx); + VARIANT2_SHUFFLE(l, idx, ax0, bx0, bx1, cx, (VARIANT == xmrig::VARIANT_RWZ ? 1 : 0)); _mm_store_si128((__m128i *)mem_out, _mm_xor_si128(bx0, cx)); } else { __m128i tmp = _mm_xor_si128(bx0, cx); @@ -558,9 +558,9 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si if (BASE == xmrig::VARIANT_2) { if (VARIANT == xmrig::VARIANT_4) { - VARIANT2_SHUFFLE(l0, idx0 & MASK, ax0, bx0, bx1, cx); + VARIANT2_SHUFFLE(l0, idx0 & MASK, ax0, bx0, bx1, cx, 0); } else { - VARIANT2_SHUFFLE2(l0, idx0 & MASK, ax0, bx0, bx1, hi, lo); + VARIANT2_SHUFFLE2(l0, idx0 & MASK, ax0, bx0, bx1, hi, lo, (VARIANT == xmrig::VARIANT_RWZ ? 1 : 0)); } } @@ -896,9 +896,9 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si if (BASE == xmrig::VARIANT_2) { if (VARIANT == xmrig::VARIANT_4) { - VARIANT2_SHUFFLE(l0, idx0 & MASK, ax0, bx00, bx01, cx0); + VARIANT2_SHUFFLE(l0, idx0 & MASK, ax0, bx00, bx01, cx0, 0); } else { - VARIANT2_SHUFFLE2(l0, idx0 & MASK, ax0, bx00, bx01, hi, lo); + VARIANT2_SHUFFLE2(l0, idx0 & MASK, ax0, bx00, bx01, hi, lo, (VARIANT == xmrig::VARIANT_RWZ ? 1 : 0)); } } @@ -952,9 +952,9 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si if (BASE == xmrig::VARIANT_2) { if (VARIANT == xmrig::VARIANT_4) { - VARIANT2_SHUFFLE(l1, idx1 & MASK, ax1, bx10, bx11, cx1); + VARIANT2_SHUFFLE(l1, idx1 & MASK, ax1, bx10, bx11, cx1, 0); } else { - VARIANT2_SHUFFLE2(l1, idx1 & MASK, ax1, bx10, bx11, hi, lo); + VARIANT2_SHUFFLE2(l1, idx1 & MASK, ax1, bx10, bx11, hi, lo, (VARIANT == xmrig::VARIANT_RWZ ? 1 : 0)); } } @@ -1056,9 +1056,9 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si lo = __umul128(idx, cl##part, &hi); \ if (BASE == xmrig::VARIANT_2) { \ if (VARIANT == xmrig::VARIANT_4) { \ - VARIANT2_SHUFFLE(l, idx & MASK, a, b0, b1, c); \ + VARIANT2_SHUFFLE(l, idx & MASK, a, b0, b1, c, 0); \ } else { \ - VARIANT2_SHUFFLE2(l, idx & MASK, a, b0, b1, hi, lo); \ + VARIANT2_SHUFFLE2(l, idx & MASK, a, b0, b1, hi, lo, (VARIANT == xmrig::VARIANT_RWZ ? 1 : 0)); \ } \ } \ if (VARIANT == xmrig::VARIANT_4) { \ diff --git a/src/workers/CpuThread.cpp b/src/workers/CpuThread.cpp index 6065bc3a..b482d0f7 100644 --- a/src/workers/CpuThread.cpp +++ b/src/workers/CpuThread.cpp @@ -323,6 +323,28 @@ xmrig::CpuThread::cn_hash_fun xmrig::CpuThread::fn(Algo algorithm, AlgoVariant a cryptonight_quad_hash, cryptonight_penta_hash, + cryptonight_single_hash, + cryptonight_double_hash, + cryptonight_single_hash, + cryptonight_double_hash, + cryptonight_triple_hash, + cryptonight_quad_hash, + cryptonight_penta_hash, + cryptonight_triple_hash, + cryptonight_quad_hash, + cryptonight_penta_hash, + + cryptonight_single_hash, + cryptonight_double_hash, + cryptonight_single_hash, + cryptonight_double_hash, + cryptonight_triple_hash, + cryptonight_quad_hash, + cryptonight_penta_hash, + cryptonight_triple_hash, + cryptonight_quad_hash, + cryptonight_penta_hash, + # ifndef XMRIG_NO_AEON cryptonight_single_hash, cryptonight_double_hash, @@ -358,6 +380,8 @@ xmrig::CpuThread::cn_hash_fun xmrig::CpuThread::fn(Algo algorithm, AlgoVariant a nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_GPU nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_WOW nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_4 + nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_WALTZ + nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_RWZ # else nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_0 nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_1 @@ -373,6 +397,8 @@ xmrig::CpuThread::cn_hash_fun xmrig::CpuThread::fn(Algo algorithm, AlgoVariant a nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_GPU nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_WOW nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_4 + nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_WALTZ + nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_RWZ # endif # ifndef XMRIG_NO_SUMO @@ -422,6 +448,8 @@ xmrig::CpuThread::cn_hash_fun xmrig::CpuThread::fn(Algo algorithm, AlgoVariant a nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_GPU nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_WOW nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_4 + nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_WALTZ + nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_RWZ # else nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_0 nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_1 @@ -437,6 +465,8 @@ xmrig::CpuThread::cn_hash_fun xmrig::CpuThread::fn(Algo algorithm, AlgoVariant a nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_GPU nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_WOW nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_4 + nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_WALTZ + nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_RWZ # endif # ifndef XMRIG_NO_CN_PICO @@ -465,6 +495,8 @@ xmrig::CpuThread::cn_hash_fun xmrig::CpuThread::fn(Algo algorithm, AlgoVariant a nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_GPU nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_WOW nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_4 + nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_WALTZ + nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_RWZ # else nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_0 nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_1 @@ -480,6 +512,8 @@ xmrig::CpuThread::cn_hash_fun xmrig::CpuThread::fn(Algo algorithm, AlgoVariant a nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_GPU nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_WOW nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_4 + nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_WALTZ + nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, // VARIANT_RWZ # endif }; diff --git a/src/workers/MultiWorker.cpp b/src/workers/MultiWorker.cpp index 340b6620..803f9e79 100644 --- a/src/workers/MultiWorker.cpp +++ b/src/workers/MultiWorker.cpp @@ -65,14 +65,16 @@ bool MultiWorker::selfTest() return false; } - const bool rc = verify(VARIANT_0, test_output_v0) && - verify(VARIANT_1, test_output_v1) && - verify(VARIANT_2, test_output_v2) && - verify(VARIANT_XTL, test_output_xtl) && - verify(VARIANT_MSR, test_output_msr) && - verify(VARIANT_XAO, test_output_xao) && - verify(VARIANT_RTO, test_output_rto) && - verify(VARIANT_HALF, test_output_half); + const bool rc = verify(VARIANT_0, test_output_v0) && + verify(VARIANT_1, test_output_v1) && + verify(VARIANT_2, test_output_v2) && + verify(VARIANT_XTL, test_output_xtl) && + verify(VARIANT_MSR, test_output_msr) && + verify(VARIANT_XAO, test_output_xao) && + verify(VARIANT_RTO, test_output_rto) && + verify(VARIANT_HALF, test_output_half) && + verify(VARIANT_WALTZ, test_output_waltz) && + verify(VARIANT_RWZ, test_output_rwz); # ifndef XMRIG_NO_CN_GPU if (!rc || N > 1) {