Cache Hash support

This commit is contained in:
En4orcer 2020-09-27 20:21:45 +01:00
parent 36ed0b4309
commit ab0876faff
10 changed files with 42 additions and 21 deletions

View file

@ -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

View file

@ -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;

View file

@ -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

View file

@ -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).

View file

@ -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;
}

View file

@ -44,7 +44,8 @@ public:
DERO,
KEVA,
RAVEN,
CONCEAL
CONCEAL,
CACHE
};

View file

@ -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<Algorithm::CN_LITE_0>::base() c
template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_HEAVY_0>::base() const { return Algorithm::CN_0; }
template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_HEAVY_XHV>::base() const { return Algorithm::CN_0; }
template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_CCX>::base() const { return Algorithm::CN_0; }
template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_CACHE_HASH>::base() const { return Algorithm::CN_0; }
template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_1>::base() const { return Algorithm::CN_1; }
template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_FAST>::base() const { return Algorithm::CN_1; }
template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_RTO>::base() const { return Algorithm::CN_1; }
@ -192,6 +200,7 @@ template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_ZLS>::iterations() con
template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_PICO_0>::iterations() const { return CN_ITER / 8; }
template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_PICO_TLO>::iterations() const { return CN_ITER / 8; }
template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_CCX>::iterations() const { return CN_ITER / 2; }
template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_CACHE_HASH>::iterations() const { return 0x50000; }
template<> constexpr inline size_t CnAlgo<Algorithm::CN_LITE_0>::memory() const { return CN_MEMORY / 2; }
@ -204,6 +213,7 @@ template<> constexpr inline size_t CnAlgo<Algorithm::CN_PICO_TLO>::memory() cons
template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_PICO_0>::mask() const { return 0x1FFF0; }
template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_CACHE_HASH>::mask() const { return 0x1FFFA0; }
} /* namespace xmrig */

View file

@ -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<Algorithm::AR2_CHUKWA>;

View file

@ -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<const __m128i *>(&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<const __m128i*>(&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);

View file

@ -612,7 +612,7 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
__m128i bx1 = _mm_set_epi64x(static_cast<int64_t>(h0[9] ^ h0[11]), static_cast<int64_t>(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<const __m128i *>(&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<const __m128i*>(&l0[idx0 & MASK]));
cryptonight_conceal_tweak(cx, conc_var);
cx = soft_aesenc(&cx, ax0, reinterpret_cast<const uint32_t*>(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<const __m128i *>(&l0[idx0 & MASK]));
cx1 = _mm_load_si128(reinterpret_cast<const __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);
}
@ -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<const __m128i*>(&l0[idx0 & MASK]));
cx1 = _mm_load_si128(reinterpret_cast<const __m128i*>(&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();
}