From 80864aef612164d3e4b86dcb8fcbd40ee787ddcd Mon Sep 17 00:00:00 2001 From: XMRig Date: Sun, 15 Nov 2020 08:24:07 +0700 Subject: [PATCH 1/5] v6.5.1-dev --- src/version.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/version.h b/src/version.h index 65e1f91..c7d81b6 100644 --- a/src/version.h +++ b/src/version.h @@ -28,14 +28,14 @@ #define APP_ID "xmrig-cuda" #define APP_NAME "XMRig" #define APP_DESC "XMRig CUDA plugin" -#define APP_VERSION "6.5.0" +#define APP_VERSION "6.5.1-dev" #define APP_DOMAIN "xmrig.com" #define APP_SITE "www.xmrig.com" #define APP_COPYRIGHT "Copyright (C) 2016-2020 xmrig.com" #define APP_VER_MAJOR 6 #define APP_VER_MINOR 5 -#define APP_VER_PATCH 0 +#define APP_VER_PATCH 1 #define API_VERSION 3 From 57716410be0ce136fffce4edad5960e673e386b2 Mon Sep 17 00:00:00 2001 From: SChernykh Date: Sat, 17 Apr 2021 14:54:04 +0200 Subject: [PATCH 2/5] Added support for Uplexa (cn/upx2 algorithm) --- CMakeLists.txt | 5 +++++ src/crypto/cn/CnAlgo.h | 21 +++++++++++++++++++++ src/crypto/common/Algorithm.cpp | 6 ++++++ src/crypto/common/Algorithm.h | 10 +++++++++- src/cuda_core.cu | 18 ++++++++++++++---- 5 files changed, 55 insertions(+), 5 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 41724d7..6e87900 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -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) @@ -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() diff --git a/src/crypto/cn/CnAlgo.h b/src/crypto/cn/CnAlgo.h index 3f59db3..1360bf6 100644 --- a/src/crypto/cn/CnAlgo.h +++ b/src/crypto/cn/CnAlgo.h @@ -66,6 +66,9 @@ class CnAlgo case Algorithm::CN_PICO: return CN_MEMORY / 8; + case Algorithm::CN_FEMTO: + return CN_MEMORY / 16; + default: break; } @@ -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; } @@ -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; } @@ -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; @@ -191,6 +208,7 @@ template<> constexpr inline Algorithm::Id CnAlgo::base() con 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; } +template<> constexpr inline Algorithm::Id CnAlgo::base() const { return Algorithm::CN_2; } template<> constexpr inline uint32_t CnAlgo::iterations() const { return CN_ITER / 2; } @@ -207,6 +225,7 @@ template<> constexpr inline uint32_t CnAlgo::iterations() con template<> constexpr inline uint32_t CnAlgo::iterations() const { return 0x60000; } 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 / 32; } template<> constexpr inline size_t CnAlgo::memory() const { return CN_MEMORY / 2; } @@ -216,9 +235,11 @@ template<> constexpr inline size_t CnAlgo::memory() co template<> constexpr inline size_t CnAlgo::memory() const { return CN_MEMORY * 2; } template<> constexpr inline size_t CnAlgo::memory() const { return CN_MEMORY / 8; } template<> constexpr inline size_t CnAlgo::memory() const { return CN_MEMORY / 8; } +template<> constexpr inline size_t CnAlgo::memory() const { return CN_MEMORY / 16; } template<> constexpr inline uint32_t CnAlgo::mask() const { return 0x1FFF0; } +template<> constexpr inline uint32_t CnAlgo::mask() const { return 0x1FFF0; } } /* namespace xmrig_cuda */ diff --git a/src/crypto/common/Algorithm.cpp b/src/crypto/common/Algorithm.cpp index 19629d0..1767ebe 100644 --- a/src/crypto/common/Algorithm.cpp +++ b/src/crypto/common/Algorithm.cpp @@ -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 diff --git a/src/crypto/common/Algorithm.h b/src/crypto/common/Algorithm.h index 37b92b8..4446df0 100644 --- a/src/crypto/common/Algorithm.h +++ b/src/crypto/common/Algorithm.h @@ -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). @@ -77,6 +78,7 @@ class Algorithm CN_LITE, CN_HEAVY, CN_PICO, + CN_FEMTO, RANDOM_X, ARGON2, ASTROBWT, @@ -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) : 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); } @@ -141,6 +143,9 @@ class Algorithm case CN_PICO: return oneMiB / 4; + case CN_FEMTO: + return oneMiB / 8; + default: break; } @@ -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: diff --git a/src/cuda_core.cu b/src/cuda_core.cu index 2d5fe77..619572b 100644 --- a/src/cuda_core.cu +++ b/src/cuda_core.cu @@ -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; } @@ -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; @@ -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(ctx, startNonce); + break; + + default: + break; + } + } } From 91187295d6fb99c6ed8a07270e35ffd7ee56b016 Mon Sep 17 00:00:00 2001 From: XMRig Date: Mon, 19 Apr 2021 20:21:03 +0700 Subject: [PATCH 3/5] v6.12.0-dev --- src/crypto/common/Algorithm.cpp | 9 +++------ src/crypto/common/Algorithm.h | 4 ++-- src/version.h | 6 +++--- 3 files changed, 8 insertions(+), 11 deletions(-) diff --git a/src/crypto/common/Algorithm.cpp b/src/crypto/common/Algorithm.cpp index 1767ebe..9d83296 100644 --- a/src/crypto/common/Algorithm.cpp +++ b/src/crypto/common/Algorithm.cpp @@ -6,8 +6,8 @@ * Copyright 2016 Jay D Dee * Copyright 2017-2018 XMR-Stak , * Copyright 2018 Lee Clagett - * Copyright 2018-2020 SChernykh - * Copyright 2016-2020 XMRig , + * Copyright 2018-2021 SChernykh + * Copyright 2016-2021 XMRig , * * This program is free software: you can redistribute it and/or modify * it under the terms of the GNU General Public License as published by @@ -82,10 +82,7 @@ static AlgoName const algorithm_names[] = { { "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 }, + { "cn/upx2", Algorithm::CN_UPX2 }, # endif # ifdef XMRIG_ALGO_ASTROBWT { "astrobwt", Algorithm::ASTROBWT_DERO }, diff --git a/src/crypto/common/Algorithm.h b/src/crypto/common/Algorithm.h index 4446df0..94dc614 100644 --- a/src/crypto/common/Algorithm.h +++ b/src/crypto/common/Algorithm.h @@ -6,8 +6,8 @@ * Copyright 2016 Jay D Dee * Copyright 2017-2018 XMR-Stak , * Copyright 2018 Lee Clagett - * Copyright 2018-2020 SChernykh - * Copyright 2016-2020 XMRig , + * Copyright 2018-2021 SChernykh + * Copyright 2016-2021 XMRig , * * This program is free software: you can redistribute it and/or modify * it under the terms of the GNU General Public License as published by diff --git a/src/version.h b/src/version.h index c7d81b6..b93942b 100644 --- a/src/version.h +++ b/src/version.h @@ -28,14 +28,14 @@ #define APP_ID "xmrig-cuda" #define APP_NAME "XMRig" #define APP_DESC "XMRig CUDA plugin" -#define APP_VERSION "6.5.1-dev" +#define APP_VERSION "6.12.0-dev" #define APP_DOMAIN "xmrig.com" #define APP_SITE "www.xmrig.com" #define APP_COPYRIGHT "Copyright (C) 2016-2020 xmrig.com" #define APP_VER_MAJOR 6 -#define APP_VER_MINOR 5 -#define APP_VER_PATCH 1 +#define APP_VER_MINOR 12 +#define APP_VER_PATCH 0 #define API_VERSION 3 From 29b21778e1f19883c955c4f3a14bb96b2ed809f7 Mon Sep 17 00:00:00 2001 From: xmrig Date: Mon, 19 Apr 2021 21:13:19 +0700 Subject: [PATCH 4/5] Update CHANGELOG.md --- CHANGELOG.md | 3 +++ 1 file changed, 3 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 00d3048..e56aa02 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -1,3 +1,6 @@ +# v6.12.0 +- [#95](https://github.com/xmrig/xmrig-cuda/pull/95) Added support for Uplexa (`cn/upx2` algorithm). + # v6.5.0 - [#74](https://github.com/xmrig/xmrig-cuda/pull/74) Fixed CUDA 8.0 support, RandomX, AstroBWT, and KawPow disabled for this CUDA version. - [#76](https://github.com/xmrig/xmrig-cuda/pull/76) Fixed high CPU usage on Cryptonight and AstroBWT. From 37d72c8b75f19fc9e4e20a7ba561978c43405f2b Mon Sep 17 00:00:00 2001 From: XMRig Date: Tue, 20 Apr 2021 19:30:02 +0700 Subject: [PATCH 5/5] v6.12.0 --- src/version.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/src/version.h b/src/version.h index b93942b..b58efe1 100644 --- a/src/version.h +++ b/src/version.h @@ -5,8 +5,8 @@ * Copyright 2014-2016 Wolf9466 * Copyright 2016 Jay D Dee * Copyright 2017-2018 XMR-Stak , - * Copyright 2018-2020 SChernykh - * Copyright 2016-2020 XMRig , + * Copyright 2018-2021 SChernykh + * Copyright 2016-2021 XMRig , * * This program is free software: you can redistribute it and/or modify * it under the terms of the GNU General Public License as published by @@ -28,7 +28,7 @@ #define APP_ID "xmrig-cuda" #define APP_NAME "XMRig" #define APP_DESC "XMRig CUDA plugin" -#define APP_VERSION "6.12.0-dev" +#define APP_VERSION "6.12.0" #define APP_DOMAIN "xmrig.com" #define APP_SITE "www.xmrig.com" #define APP_COPYRIGHT "Copyright (C) 2016-2020 xmrig.com"