Skip to content

Commit

Permalink
Added support for Uplexa (cn/upx2 algorithm)
Browse files Browse the repository at this point in the history
  • Loading branch information
@SChernykh
SChernykh committed Apr 17, 2021
1 parent 80864ae commit 5771641
Show file tree
Hide file tree
Showing 5 changed files with 55 additions and 5 deletions.
5 changes: 5 additions & 0 deletions CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -10,6 +10,7 @@ option(WITH_CN_R "Enable CryptoNight-R algorithm" ON)
option(WITH_CN_LITE "Enable CryptoNight-Lite algorithms family" ON)
option(WITH_CN_HEAVY "Enable CryptoNight-Heavy algorithms family" ON)
option(WITH_CN_PICO "Enable CryptoNight-Pico algorithm" ON)
option(WITH_CN_FEMTO "Enable CryptoNight-UPX2 algorithm" ON)
option(WITH_ARGON2 "Enable Argon2 algorithms family" OFF) #unsupported

if (CUDA_VERSION VERSION_LESS 9.0)
Expand All @@ -35,6 +36,10 @@ if (WITH_CN_PICO)
add_definitions(/DXMRIG_ALGO_CN_PICO)
endif()

if (WITH_CN_FEMTO)
add_definitions(/DXMRIG_ALGO_CN_FEMTO)
endif()

if (WITH_RANDOMX)
add_definitions(/DXMRIG_ALGO_RANDOMX)
endif()
Expand Down
21 changes: 21 additions & 0 deletions src/crypto/cn/CnAlgo.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -66,6 +66,9 @@ class CnAlgo
case Algorithm::CN_PICO:
return CN_MEMORY / 8;

case Algorithm::CN_FEMTO:
return CN_MEMORY / 16;

default:
break;
}
Expand Down Expand Up @@ -111,6 +114,11 @@ class CnAlgo
return CN_ITER / 8;
# endif

# ifdef XMRIG_ALGO_CN_FEMTO
case Algorithm::CN_UPX2:
return CN_ITER / 32;
# endif

default:
break;
}
Expand All @@ -126,6 +134,12 @@ class CnAlgo
}
# endif

# ifdef XMRIG_ALGO_CN_FEMTO
if (algo == Algorithm::CN_UPX2) {
return 0x1FFF0;
}
# endif

return ((memory(algo) - 1) / 16) * 16;
}

Expand Down Expand Up @@ -164,6 +178,9 @@ class CnAlgo
# ifdef XMRIG_ALGO_CN_PICO
case Algorithm::CN_PICO_0:
case Algorithm::CN_PICO_TLO:
# endif
# ifdef XMRIG_ALGO_CN_FEMTO
case Algorithm::CN_UPX2:
# endif
return Algorithm::CN_2;

Expand Down Expand Up @@ -191,6 +208,7 @@ template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_FAST>::base() con
template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_RTO>::base() const { return Algorithm::CN_1; }
template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_LITE_1>::base() const { return Algorithm::CN_1; }
template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_HEAVY_TUBE>::base() const { return Algorithm::CN_1; }
template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_UPX2>::base() const { return Algorithm::CN_2; }


template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_FAST>::iterations() const { return CN_ITER / 2; }
Expand All @@ -207,6 +225,7 @@ template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_RWZ>::iterations() con
template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_ZLS>::iterations() const { return 0x60000; }
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_UPX2>::iterations() const { return CN_ITER / 32; }


template<> constexpr inline size_t CnAlgo<Algorithm::CN_LITE_0>::memory() const { return CN_MEMORY / 2; }
Expand All @@ -216,9 +235,11 @@ template<> constexpr inline size_t CnAlgo<Algorithm::CN_HEAVY_TUBE>::memory() co
template<> constexpr inline size_t CnAlgo<Algorithm::CN_HEAVY_XHV>::memory() const { return CN_MEMORY * 2; }
template<> constexpr inline size_t CnAlgo<Algorithm::CN_PICO_0>::memory() const { return CN_MEMORY / 8; }
template<> constexpr inline size_t CnAlgo<Algorithm::CN_PICO_TLO>::memory() const { return CN_MEMORY / 8; }
template<> constexpr inline size_t CnAlgo<Algorithm::CN_UPX2>::memory() const { return CN_MEMORY / 16; }


template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_PICO_0>::mask() const { return 0x1FFF0; }
template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_UPX2>::mask() const { return 0x1FFF0; }


} /* namespace xmrig_cuda */
Expand Down
6 changes: 6 additions & 0 deletions src/crypto/common/Algorithm.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -81,6 +81,12 @@ static AlgoName const algorithm_names[] = {
{ "cn-pico", Algorithm::CN_PICO_0 },
{ "cn-pico/tlo", Algorithm::CN_PICO_TLO },
# endif
# ifdef XMRIG_ALGO_CN_FEMTO
{ "cn/upx2", Algorithm::CN_UPX2 },
// Algo names from other miners
{ "cn-extremelite/upx2", Algorithm::CN_UPX2 },
{ "cryptonight-upx/2", Algorithm::CN_UPX2 },
# endif
# ifdef XMRIG_ALGO_ASTROBWT
{ "astrobwt", Algorithm::ASTROBWT_DERO },
# endif
Expand Down
10 changes: 9 additions & 1 deletion src/crypto/common/Algorithm.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -58,6 +58,7 @@ class Algorithm
CN_PICO_0, // "cn-pico" CryptoNight-Pico
CN_PICO_TLO, // "cn-pico/tlo" CryptoNight-Pico (TLO)
CN_CCX, // "cn/ccx" Conceal (CCX)
CN_UPX2, // "cn/upx2" Uplexa (UPX2)
RX_0, // "rx/0" RandomX (reference configuration).
RX_WOW, // "rx/wow" RandomWOW (Wownero).
RX_ARQ, // "rx/arq" RandomARQ (Arqma).
Expand All @@ -77,6 +78,7 @@ class Algorithm
CN_LITE,
CN_HEAVY,
CN_PICO,
CN_FEMTO,
RANDOM_X,
ARGON2,
ASTROBWT,
Expand All @@ -88,7 +90,7 @@ class Algorithm
inline Algorithm(Id id) : m_id(id) {}
inline Algorithm(int id) : m_id(id > INVALID && id < MAX ? static_cast<Id>(id) : INVALID) {}

inline bool isCN() const { auto f = family(); return f == CN || f == CN_LITE || f == CN_HEAVY || f == CN_PICO; }
inline bool isCN() const { auto f = family(); return f == CN || f == CN_LITE || f == CN_HEAVY || f == CN_PICO || f == CN_FEMTO; }
inline bool isEqual(const Algorithm &other) const { return m_id == other.m_id; }
inline bool isValid() const { return m_id != INVALID; }
inline Family family() const { return family(m_id); }
Expand Down Expand Up @@ -141,6 +143,9 @@ class Algorithm
case CN_PICO:
return oneMiB / 4;

case CN_FEMTO:
return oneMiB / 8;

default:
break;
}
Expand Down Expand Up @@ -229,6 +234,9 @@ class Algorithm
case CN_PICO_TLO:
return CN_PICO;

case CN_UPX2:
return CN_FEMTO;

case RX_0:
case RX_WOW:
case RX_ARQ:
Expand Down
18 changes: 14 additions & 4 deletions src/cuda_core.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -350,8 +350,8 @@ __global__ void cryptonight_core_gpu_phase2_double(
__syncthreads();
# endif

myChunks[idx1 ^ 2 + sub] = ((ALGO == Algorithm::CN_RWZ) ? chunk1 : chunk3) + bx1;
myChunks[idx1 ^ 4 + sub] = ((ALGO == Algorithm::CN_RWZ) ? chunk3 : chunk1) + bx0;
myChunks[idx1 ^ 2 + sub] = (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? chunk1 : chunk3) + bx1;
myChunks[idx1 ^ 4 + sub] = (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? chunk3 : chunk1) + bx0;
myChunks[idx1 ^ 6 + sub] = chunk2 + ax0;
}

Expand Down Expand Up @@ -405,8 +405,8 @@ __global__ void cryptonight_core_gpu_phase2_double(
__syncthreads( );
# endif

myChunks[idx1 ^ 2 + sub] = ((ALGO == Algorithm::CN_RWZ) ? chunk1 : chunk3) + bx1;
myChunks[idx1 ^ 4 + sub] = ((ALGO == Algorithm::CN_RWZ) ? chunk3 : chunk1) + bx0;
myChunks[idx1 ^ 2 + sub] = (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? chunk1 : chunk3) + bx1;
myChunks[idx1 ^ 4 + sub] = (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? chunk3 : chunk1) + bx0;
myChunks[idx1 ^ 6 + sub] = chunk2 + ax0;

ax0 += res;
Expand Down Expand Up @@ -945,4 +945,14 @@ void cryptonight_gpu_hash(nvid_ctx *ctx, const xmrig_cuda::Algorithm &algorithm,
break;
}
}
else if (algorithm.family() == Algorithm::CN_FEMTO) {
switch (algorithm.id()) {
case Algorithm::CN_UPX2:
cryptonight_core_gpu_hash<Algorithm::CN_UPX2>(ctx, startNonce);
break;

default:
break;
}
}
}

0 comments on commit 5771641

Please sign in to comment.