diff --git a/src/backend/opencl/cl/cn/algorithm.cl b/src/backend/opencl/cl/cn/algorithm.cl index 72ba9936..49e55775 100644 --- a/src/backend/opencl/cl/cn/algorithm.cl +++ b/src/backend/opencl/cl/cn/algorithm.cl @@ -27,6 +27,7 @@ #define ALGO_AR2_WRKZ 26 #define ALGO_ASTROBWT_DERO 27 #define ALGO_KAWPOW_RVN 28 +#define ALGO_CN_CACHE_HASH 29 #define FAMILY_UNKNOWN 0 #define FAMILY_CN 1 diff --git a/src/backend/opencl/cl/cn/cryptonight.cl b/src/backend/opencl/cl/cn/cryptonight.cl index 8b69185b..ea037311 100644 --- a/src/backend/opencl/cl/cn/cryptonight.cl +++ b/src/backend/opencl/cl/cn/cryptonight.cl @@ -244,7 +244,7 @@ __kernel void cn1(__global ulong *input, __global uint4 *Scratchpad, __global ul { uint idx0 = a[0]; -# if (ALGO == ALGO_CN_CCX) +# if ((ALGO == ALGO_CN_CCX) || (ALGO == ALGO_CN_CACHE_HASH)) float4 conc_var = (float4)(0.0f); const uint4 conc_t = (uint4)(0x807FFFFFU); const uint4 conc_u = (uint4)(0x40000000U); @@ -257,7 +257,7 @@ __kernel void cn1(__global ulong *input, __global uint4 *Scratchpad, __global ul ((uint4 *)c)[0] = Scratchpad[IDX((idx0 & MASK) >> 4)]; -# if (ALGO == ALGO_CN_CCX) +# if ((ALGO == ALGO_CN_CCX) || (ALGO == ALGO_CN_CACHE_HASH)) { float4 r = convert_float4_rte(((int4 *)c)[0]) + conc_var; r = r * r * r; diff --git a/src/base/crypto/Algorithm.cpp b/src/base/crypto/Algorithm.cpp index 778bee76..b7898e04 100644 --- a/src/base/crypto/Algorithm.cpp +++ b/src/base/crypto/Algorithm.cpp @@ -129,6 +129,8 @@ static AlgoName const algorithm_names[] = { # endif { "cryptonight/ccx", "cn/ccx", Algorithm::CN_CCX }, { "cryptonight/conceal", "cn/conceal", Algorithm::CN_CCX }, + { "cryptonight/cache_hash", "cn/cache_hash", Algorithm::CN_CACHE_HASH }, + { "cryptonight/cache", "cn/cache", Algorithm::CN_CACHE_HASH }, }; @@ -307,6 +309,7 @@ xmrig::Algorithm::Family xmrig::Algorithm::family(Id id) case CN_ZLS: case CN_DOUBLE: case CN_CCX: + case CN_CACHE_HASH: return CN; # ifdef XMRIG_ALGO_CN_LITE diff --git a/src/base/crypto/Algorithm.h b/src/base/crypto/Algorithm.h index b04765d5..335abe39 100644 --- a/src/base/crypto/Algorithm.h +++ b/src/base/crypto/Algorithm.h @@ -64,6 +64,7 @@ public: CN_PICO_0, // "cn-pico" CryptoNight-Pico CN_PICO_TLO, // "cn-pico/tlo" CryptoNight-Pico (TLO) CN_CCX, // "cn/ccx" Conceal (CCX) + CN_CACHE_HASH, // "cn/cache_hash" Cache (CXCHE) RX_0, // "rx/0" RandomX (reference configuration). RX_WOW, // "rx/wow" RandomWOW (Wownero). RX_LOKI, // "rx/loki" RandomXL (Loki). diff --git a/src/base/crypto/Coin.cpp b/src/base/crypto/Coin.cpp index 33115484..ba379bdd 100644 --- a/src/base/crypto/Coin.cpp +++ b/src/base/crypto/Coin.cpp @@ -56,7 +56,8 @@ static CoinName const coin_names[] = { { "ravencoin", Coin::RAVEN }, { "raven", Coin::RAVEN }, { "rvn", Coin::RAVEN }, - { "conceal", Coin::CONCEAL } + { "conceal", Coin::CONCEAL }, + { "cache", Coin::CACHE } }; @@ -85,6 +86,9 @@ xmrig::Algorithm::Id xmrig::Coin::algorithm(uint8_t blobVersion) const case CONCEAL: return Algorithm::CN_CCX; + case CACHE: + return Algorithm::CN_CACHE_HASH; + case INVALID: break; } diff --git a/src/base/crypto/Coin.h b/src/base/crypto/Coin.h index 084672fc..03c5c9c7 100644 --- a/src/base/crypto/Coin.h +++ b/src/base/crypto/Coin.h @@ -44,7 +44,8 @@ public: DERO, KEVA, RAVEN, - CONCEAL + CONCEAL, + CACHE }; diff --git a/src/crypto/cn/CnAlgo.h b/src/crypto/cn/CnAlgo.h index 5bca1283..b79b8bc4 100644 --- a/src/crypto/cn/CnAlgo.h +++ b/src/crypto/cn/CnAlgo.h @@ -81,6 +81,8 @@ public: # endif case Algorithm::CN_CCX: return CN_ITER / 2; + case Algorithm::CN_CACHE_HASH: + return 0x50000; case Algorithm::CN_RWZ: case Algorithm::CN_ZLS: @@ -111,6 +113,10 @@ public: } # endif + if (algo == Algorithm::CN_CACHE_HASH){ + return 0x1FFFA0; + } + return ((memory(algo) - 1) / 16) * 16; } @@ -127,6 +133,7 @@ public: case Algorithm::CN_HEAVY_XHV: # endif case Algorithm::CN_CCX: + case Algorithm::CN_CACHE_HASH: return Algorithm::CN_0; case Algorithm::CN_1: @@ -171,6 +178,7 @@ template<> constexpr inline Algorithm::Id CnAlgo::base() c template<> constexpr inline Algorithm::Id CnAlgo::base() const { return Algorithm::CN_0; } template<> constexpr inline Algorithm::Id CnAlgo::base() const { return Algorithm::CN_0; } template<> constexpr inline Algorithm::Id CnAlgo::base() const { return Algorithm::CN_0; } +template<> constexpr inline Algorithm::Id CnAlgo::base() const { return Algorithm::CN_0; } template<> constexpr inline Algorithm::Id CnAlgo::base() const { return Algorithm::CN_1; } template<> constexpr inline Algorithm::Id CnAlgo::base() const { return Algorithm::CN_1; } template<> constexpr inline Algorithm::Id CnAlgo::base() const { return Algorithm::CN_1; } @@ -192,6 +200,7 @@ template<> constexpr inline uint32_t CnAlgo::iterations() con template<> constexpr inline uint32_t CnAlgo::iterations() const { return CN_ITER / 8; } template<> constexpr inline uint32_t CnAlgo::iterations() const { return CN_ITER / 8; } template<> constexpr inline uint32_t CnAlgo::iterations() const { return CN_ITER / 2; } +template<> constexpr inline uint32_t CnAlgo::iterations() const { return 0x50000; } template<> constexpr inline size_t CnAlgo::memory() const { return CN_MEMORY / 2; } @@ -204,6 +213,7 @@ template<> constexpr inline size_t CnAlgo::memory() cons template<> constexpr inline uint32_t CnAlgo::mask() const { return 0x1FFF0; } +template<> constexpr inline uint32_t CnAlgo::mask() const { return 0x1FFFA0; } } /* namespace xmrig */ diff --git a/src/crypto/cn/CnHash.cpp b/src/crypto/cn/CnHash.cpp index 9f519492..518e83e8 100644 --- a/src/crypto/cn/CnHash.cpp +++ b/src/crypto/cn/CnHash.cpp @@ -271,6 +271,7 @@ xmrig::CnHash::CnHash() # endif ADD_FN(Algorithm::CN_CCX); + ADD_FN(Algorithm::CN_CACHE_HASH); # ifdef XMRIG_ALGO_ARGON2 m_map[Algorithm::AR2_CHUKWA][AV_SINGLE][Assembly::NONE] = argon2::single_hash; diff --git a/src/crypto/cn/CryptoNight_arm.h b/src/crypto/cn/CryptoNight_arm.h index 9899c27a..2e62a037 100644 --- a/src/crypto/cn/CryptoNight_arm.h +++ b/src/crypto/cn/CryptoNight_arm.h @@ -465,7 +465,7 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si __m128i bx1 = _mm_set_epi64x(h0[9] ^ h0[11], h0[8] ^ h0[10]); __m128 conc_var; - if (ALGO == Algorithm::CN_CCX) { + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { conc_var = _mm_setzero_ps(); } @@ -475,7 +475,7 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si __m128i cx; if (IS_CN_HEAVY_TUBE || !SOFT_AES) { cx = _mm_load_si128(reinterpret_cast(&l0[idx0 & MASK])); - if (ALGO == Algorithm::CN_CCX) { + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { cryptonight_conceal_tweak(cx, conc_var); } } @@ -485,7 +485,7 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si cx = aes_round_tweak_div(cx, ax0); } else if (SOFT_AES) { - if (ALGO == Algorithm::CN_CCX) { + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { cx = _mm_load_si128(reinterpret_cast(&l0[idx0 & MASK])); cryptonight_conceal_tweak(cx, conc_var); cx = soft_aesenc((uint32_t*)&cx, ax0); @@ -627,7 +627,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si __m128i bx11 = _mm_set_epi64x(h1[9] ^ h1[11], h1[8] ^ h1[10]); __m128 conc_var0, conc_var1; - if (ALGO == Algorithm::CN_CCX) { + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { conc_var0 = _mm_setzero_ps(); conc_var1 = _mm_setzero_ps(); } @@ -640,7 +640,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si if (IS_CN_HEAVY_TUBE || !SOFT_AES) { cx0 = _mm_load_si128((__m128i *) &l0[idx0 & MASK]); cx1 = _mm_load_si128((__m128i *) &l1[idx1 & MASK]); - if (ALGO == Algorithm::CN_CCX) { + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { cryptonight_conceal_tweak(cx0, conc_var0); cryptonight_conceal_tweak(cx1, conc_var1); } @@ -653,7 +653,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si cx1 = aes_round_tweak_div(cx1, ax1); } else if (SOFT_AES) { - if (ALGO == Algorithm::CN_CCX) { + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { cx0 = _mm_load_si128((__m128i *) &l0[idx0 & MASK]); cx1 = _mm_load_si128((__m128i *) &l1[idx1 & MASK]); cryptonight_conceal_tweak(cx0, conc_var0); diff --git a/src/crypto/cn/CryptoNight_x86.h b/src/crypto/cn/CryptoNight_x86.h index 48748cf0..4d75cae6 100644 --- a/src/crypto/cn/CryptoNight_x86.h +++ b/src/crypto/cn/CryptoNight_x86.h @@ -612,7 +612,7 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si __m128i bx1 = _mm_set_epi64x(static_cast(h0[9] ^ h0[11]), static_cast(h0[8] ^ h0[10])); __m128 conc_var; - if (ALGO == Algorithm::CN_CCX) { + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { conc_var = _mm_setzero_ps(); RESTORE_ROUNDING_MODE(); } @@ -621,7 +621,7 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si __m128i cx; if (IS_CN_HEAVY_TUBE || !SOFT_AES) { cx = _mm_load_si128(reinterpret_cast(&l0[idx0 & MASK])); - if (ALGO == Algorithm::CN_CCX) { + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { cryptonight_conceal_tweak(cx, conc_var); } } @@ -631,7 +631,7 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si cx = aes_round_tweak_div(cx, ax0); } else if (SOFT_AES) { - if (ALGO == Algorithm::CN_CCX) { + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { cx = _mm_load_si128(reinterpret_cast(&l0[idx0 & MASK])); cryptonight_conceal_tweak(cx, conc_var); cx = soft_aesenc(&cx, ax0, reinterpret_cast(saes_table)); @@ -1005,7 +1005,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si __m128i bx11 = _mm_set_epi64x(h1[9] ^ h1[11], h1[8] ^ h1[10]); __m128 conc_var0, conc_var1; - if (ALGO == Algorithm::CN_CCX) { + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { conc_var0 = _mm_setzero_ps(); conc_var1 = _mm_setzero_ps(); RESTORE_ROUNDING_MODE(); @@ -1019,7 +1019,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si if (IS_CN_HEAVY_TUBE || !SOFT_AES) { cx0 = _mm_load_si128(reinterpret_cast(&l0[idx0 & MASK])); cx1 = _mm_load_si128(reinterpret_cast(&l1[idx1 & MASK])); - if (ALGO == Algorithm::CN_CCX) { + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { cryptonight_conceal_tweak(cx0, conc_var0); cryptonight_conceal_tweak(cx1, conc_var1); } @@ -1032,7 +1032,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si cx1 = aes_round_tweak_div(cx1, ax1); } else if (SOFT_AES) { - if (ALGO == Algorithm::CN_CCX) { + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { cx0 = _mm_load_si128(reinterpret_cast(&l0[idx0 & MASK])); cx1 = _mm_load_si128(reinterpret_cast(&l1[idx1 & MASK])); cryptonight_conceal_tweak(cx0, conc_var0); @@ -1201,7 +1201,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si #define CN_STEP1(a, b0, b1, c, l, ptr, idx, conc_var) \ ptr = reinterpret_cast<__m128i*>(&l[idx & MASK]); \ c = _mm_load_si128(ptr); \ - if (ALGO == Algorithm::CN_CCX) { \ + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { \ cryptonight_conceal_tweak(c, conc_var); \ } @@ -1305,7 +1305,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si __m128i bx##n##1 = _mm_set_epi64x(h##n[9] ^ h##n[11], h##n[8] ^ h##n[10]); \ __m128i cx##n = _mm_setzero_si128(); \ __m128 conc_var##n; \ - if (ALGO == Algorithm::CN_CCX) { \ + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { \ conc_var##n = _mm_setzero_ps(); \ } \ VARIANT4_RANDOM_MATH_INIT(n); @@ -1347,7 +1347,7 @@ inline void cryptonight_triple_hash(const uint8_t *__restrict__ input, size_t si CONST_INIT(ctx[1], 1); CONST_INIT(ctx[2], 2); VARIANT2_SET_ROUNDING_MODE(); - if (ALGO == Algorithm::CN_CCX) { + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { RESTORE_ROUNDING_MODE(); } @@ -1424,7 +1424,7 @@ inline void cryptonight_quad_hash(const uint8_t *__restrict__ input, size_t size CONST_INIT(ctx[2], 2); CONST_INIT(ctx[3], 3); VARIANT2_SET_ROUNDING_MODE(); - if (ALGO == Algorithm::CN_CCX) { + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { RESTORE_ROUNDING_MODE(); } @@ -1509,7 +1509,7 @@ inline void cryptonight_penta_hash(const uint8_t *__restrict__ input, size_t siz CONST_INIT(ctx[3], 3); CONST_INIT(ctx[4], 4); VARIANT2_SET_ROUNDING_MODE(); - if (ALGO == Algorithm::CN_CCX) { + if ((ALGO == Algorithm::CN_CCX) || (ALGO == Algorithm::CN_CACHE_HASH)) { RESTORE_ROUNDING_MODE(); }