Merge ab0876faff
into 1bcfd0cdea
This commit is contained in:
commit
577a51fd3c
10 changed files with 42 additions and 21 deletions
|
@ -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
|
||||
|
|
|
@ -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;
|
||||
|
|
|
@ -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 },
|
||||
};
|
||||
|
||||
|
||||
|
@ -308,6 +310,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
|
||||
|
|
|
@ -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_ARQ, // "rx/arq" RandomARQ (Arqma).
|
||||
|
|
|
@ -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;
|
||||
}
|
||||
|
|
|
@ -44,7 +44,8 @@ public:
|
|||
DERO,
|
||||
KEVA,
|
||||
RAVEN,
|
||||
CONCEAL
|
||||
CONCEAL,
|
||||
CACHE
|
||||
};
|
||||
|
||||
|
||||
|
|
|
@ -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 */
|
||||
|
|
|
@ -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>;
|
||||
|
|
|
@ -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);
|
||||
|
|
|
@ -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();
|
||||
}
|
||||
|
||||
|
|
Loading…
Add table
Add a link
Reference in a new issue