Merge 3bc2f27b97 into e32731b60b

Detect AMD engineering samples in randomx_boost.sh

CHANGELOG.md

@@ -1,3 +1,14 @@
+# v6.22.0
+- [#2411](https://github.com/xmrig/xmrig/pull/2411) Added support for [Yada](https://yadacoin.io/) (`rx/yada` algorithm).
+- [#3492](https://github.com/xmrig/xmrig/pull/3492) Fixed `--background` option on Unix systems.
+- [#3518](https://github.com/xmrig/xmrig/pull/3518) Possible fix for corrupted API output in rare cases.
+- [#3522](https://github.com/xmrig/xmrig/pull/3522) Removed `rx/keva` algorithm.
+- [#3525](https://github.com/xmrig/xmrig/pull/3525) Added Zen5 detection.
+- [#3528](https://github.com/xmrig/xmrig/pull/3528) Added `rx/yada` OpenCL support.
+
+# v6.21.3
+- [#3462](https://github.com/xmrig/xmrig/pull/3462) RandomX: correct memcpy size for JIT initialization.
+
 # v6.21.2
 - The dependencies of all prebuilt releases have been updated. Support for old Ubuntu releases has been dropped.
 - [#2800](https://github.com/xmrig/xmrig/issues/2800) Fixed donation with GhostRider algorithm for builds without KawPow algorithm.

doc/ALGORITHMS.md

@@ -13,7 +13,6 @@ Option `coin` useful for pools without [algorithm negotiation](https://xmrig.com
 | Name | Memory | Version | Description | Notes |
 |------|--------|---------|-------------|-------|
 | `kawpow` | - | 6.0.0+ | KawPow (Ravencoin) | GPU only |
-| `rx/keva` | 1 MB | 5.9.0+ | RandomKEVA (RandomX variant for Keva). |  |
 | `astrobwt` | 20 MB | 5.8.0+ | AstroBWT (Dero). |  |
 | `cn-pico/tlo` | 256 KB | 5.5.0+ | CryptoNight-Pico (Talleo). |  |
 | `rx/sfx` | 2 MB | 5.4.0+ | RandomSFX (RandomX variant for Safex). |  |

scripts/build.hwloc.sh

@@ -1,8 +1,8 @@
 #!/bin/sh -e
 
 HWLOC_VERSION_MAJOR="2"
-HWLOC_VERSION_MINOR="10"
-HWLOC_VERSION_PATCH="0"
+HWLOC_VERSION_MINOR="11"
+HWLOC_VERSION_PATCH="2"
 
 HWLOC_VERSION="${HWLOC_VERSION_MAJOR}.${HWLOC_VERSION_MINOR}.${HWLOC_VERSION_PATCH}"
 

scripts/build.openssl.sh

@@ -1,6 +1,6 @@
 #!/bin/sh -e
 
-OPENSSL_VERSION="1.1.1s"
+OPENSSL_VERSION="1.1.1u"
 
 mkdir -p deps
 mkdir -p deps/include
@@ -8,7 +8,7 @@ mkdir -p deps/lib
 
 mkdir -p build && cd build
 
-wget https://www.openssl.org/source/openssl-${OPENSSL_VERSION}.tar.gz -O openssl-${OPENSSL_VERSION}.tar.gz
+wget https://openssl.org/source/old/1.1.1/openssl-${OPENSSL_VERSION}.tar.gz -O openssl-${OPENSSL_VERSION}.tar.gz
 tar -xzf openssl-${OPENSSL_VERSION}.tar.gz
 
 cd openssl-${OPENSSL_VERSION}

scripts/build.openssl3.sh

@@ -1,6 +1,6 @@
 #!/bin/sh -e
 
-OPENSSL_VERSION="3.0.13"
+OPENSSL_VERSION="3.0.15"
 
 mkdir -p deps
 mkdir -p deps/include
@@ -8,7 +8,7 @@ mkdir -p deps/lib
 
 mkdir -p build && cd build
 
-wget https://www.openssl.org/source/openssl-${OPENSSL_VERSION}.tar.gz -O openssl-${OPENSSL_VERSION}.tar.gz
+wget https://github.com/openssl/openssl/releases/download/openssl-${OPENSSL_VERSION}/openssl-${OPENSSL_VERSION}.tar.gz -O openssl-${OPENSSL_VERSION}.tar.gz
 tar -xzf openssl-${OPENSSL_VERSION}.tar.gz
 
 cd openssl-${OPENSSL_VERSION}

scripts/build.uv.sh

@@ -1,6 +1,6 @@
 #!/bin/sh -e
 
-UV_VERSION="1.48.0"
+UV_VERSION="1.49.2"
 
 mkdir -p deps
 mkdir -p deps/include

scripts/generate_cl.js

@@ -50,7 +50,6 @@ function rx()
         'randomx_constants_monero.h',
         'randomx_constants_wow.h',
         'randomx_constants_arqma.h',
-        'randomx_constants_keva.h',
         'randomx_constants_graft.h',
         'aes.cl',
         'blake2b.cl',

scripts/randomx_boost.sh

@@ -8,7 +8,7 @@ else
 	modprobe msr allow_writes=on
 fi
 
-if grep -E 'AMD Ryzen|AMD EPYC' /proc/cpuinfo > /dev/null;
+if grep -E 'AMD Ryzen|AMD EPYC|AuthenticAMD' /proc/cpuinfo > /dev/null;
 	then
 	if grep "cpu family[[:space:]]\{1,\}:[[:space:]]25" /proc/cpuinfo > /dev/null;
 		then
@@ -28,6 +28,14 @@ if grep -E 'AMD Ryzen|AMD EPYC' /proc/cpuinfo > /dev/null;
 					wrmsr -a 0xc001102b 0x2000cc10
 					echo "MSR register values for Zen3 applied"
 				fi
+		elif grep "cpu family[[:space:]]\{1,\}:[[:space:]]26" /proc/cpuinfo > /dev/null;
+			then
+				echo "Detected Zen5 CPU"
+				wrmsr -a 0xc0011020 0x4400000000000
+				wrmsr -a 0xc0011021 0x4000000000040
+				wrmsr -a 0xc0011022 0x8680000401570000
+				wrmsr -a 0xc001102b 0x2040cc10
+				echo "MSR register values for Zen5 applied"
 		else
 			echo "Detected Zen1/Zen2 CPU"
 			wrmsr -a 0xc0011020 0

src/App.cpp

@@ -6,8 +6,8 @@
  * Copyright 2016      Jay D Dee   <jayddee246@gmail.com>
  * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
  * Copyright 2018      Lee Clagett <https://github.com/vtnerd>
- * Copyright 2018-2020 SChernykh   <https://github.com/SChernykh>
- * Copyright 2016-2020 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2024 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
  *
  *   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
@@ -23,7 +23,6 @@
  *   along with this program. If not, see <http://www.gnu.org/licenses/>.
  */
 
-
 #include <cstdlib>
 #include <uv.h>
 
@@ -61,13 +60,13 @@ int xmrig::App::exec()
         return 2;
     }
 
-    m_signals = std::make_shared<Signals>(this);
-
     int rc = 0;
     if (background(rc)) {
         return rc;
     }
 
+    m_signals = std::make_shared<Signals>(this);
+
     rc = m_controller->init();
     if (rc != 0) {
         return rc;

src/App_unix.cpp

@@ -5,8 +5,8 @@
  * Copyright 2014-2016 Wolf9466    <https://github.com/OhGodAPet>
  * Copyright 2016      Jay D Dee   <jayddee246@gmail.com>
  * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
- * Copyright 2018-2020 SChernykh   <https://github.com/SChernykh>
- * Copyright 2016-2020 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2024 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
  *
  *   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
@@ -22,7 +22,6 @@
  *   along with this program. If not, see <http://www.gnu.org/licenses/>.
  */
 
-
 #include <cstdlib>
 #include <csignal>
 #include <cerrno>
@@ -53,16 +52,9 @@ bool xmrig::App::background(int &rc)
         return true;
     }
 
-    i = setsid();
-
-    if (i < 0) {
+    if (setsid() < 0) {
         LOG_ERR("setsid() failed (errno = %d)", errno);
     }
 
-    i = chdir("/");
-    if (i < 0) {
-        LOG_ERR("chdir() failed (errno = %d)", errno);
-    }
-
     return false;
 }

src/backend/cpu/CpuConfig_gen.h

@@ -1,6 +1,6 @@
 /* XMRig
- * Copyright (c) 2018-2021 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2021 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright (c) 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright (c) 2016-2024 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
  *
  *   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
@@ -122,17 +122,6 @@ size_t inline generate<Algorithm::RANDOM_X>(Threads<CpuThreads> &threads, uint32
         }
     }
 
-    if (!threads.isExist(Algorithm::RX_KEVA)) {
-        auto keva = cpuInfo->threads(Algorithm::RX_KEVA, limit);
-        if (keva == wow) {
-            threads.setAlias(Algorithm::RX_KEVA, Algorithm::kRX_WOW);
-            ++count;
-        }
-        else {
-            count += threads.move(Algorithm::kRX_KEVA, std::move(keva));
-        }
-    }
-
     if (!threads.isExist(Algorithm::RX_WOW)) {
         count += threads.move(Algorithm::kRX_WOW, std::move(wow));
     }

src/backend/cpu/CpuWorker.cpp

@@ -67,7 +67,8 @@ VirtualMemory* cn_heavyZen3Memory = nullptr;
 template<size_t N>
 xmrig::CpuWorker<N>::CpuWorker(size_t id, const CpuLaunchData &data) :
     Worker(id, data.affinity, data.priority),
-    m_algorithm(data.algorithm),
+    m_algorithm_family(data.algorithm.family()),
+    m_algorithm_l3(data.algorithm.l3()),
     m_assembly(data.assembly),
     m_hwAES(data.hwAES),
     m_yield(data.yield),
@@ -82,24 +83,29 @@ xmrig::CpuWorker<N>::CpuWorker(size_t id, const CpuLaunchData &data) :
     const uint32_t model = Cpu::info()->model();
     const bool is_vermeer = (arch == ICpuInfo::ARCH_ZEN3) && (model == 0x21);
     const bool is_raphael = (arch == ICpuInfo::ARCH_ZEN4) && (model == 0x61);
-    if ((N == 1) && (m_av == CnHash::AV_SINGLE) && (m_algorithm.family() == Algorithm::CN_HEAVY) && (m_assembly != Assembly::NONE) && (is_vermeer || is_raphael)) {
+    if ((N == 1) && (m_av == CnHash::AV_SINGLE) && (m_algorithm_family == Algorithm::CN_HEAVY) && (m_assembly != Assembly::NONE) && (is_vermeer || is_raphael)) {
         std::lock_guard<std::mutex> lock(cn_heavyZen3MemoryMutex);
         if (!cn_heavyZen3Memory) {
             // Round up number of threads to the multiple of 8
             const size_t num_threads = ((m_threads + 7) / 8) * 8;
-            cn_heavyZen3Memory = new VirtualMemory(m_algorithm.l3() * num_threads, data.hugePages, false, false, node());
+            cn_heavyZen3Memory = new VirtualMemory(m_algorithm_l3 * num_threads, data.hugePages, false, false, node());
         }
         m_memory = cn_heavyZen3Memory;
     }
     else
 #   endif
     {
-        m_memory = new VirtualMemory(m_algorithm.l3() * N, data.hugePages, false, true, node());
+        m_memory = new VirtualMemory(m_algorithm_l3 * N, data.hugePages, false, true, node());
     }
 
 #   ifdef XMRIG_ALGO_GHOSTRIDER
     m_ghHelper = ghostrider::create_helper_thread(affinity(), data.priority, data.affinities);
 #   endif
+
+    std::vector<Algorithm> family = Algorithm::all([this](const Algorithm &algo) -> bool { return m_algorithm_family == algo.family(); });
+    for (auto& algorithm : family) {
+        m_hash_fns.emplace(algorithm.id(), CnHash::fn(algorithm, m_av, m_assembly));
+    }
 }
 
 
@@ -124,6 +130,23 @@ xmrig::CpuWorker<N>::~CpuWorker()
 #   endif
 }
 
+template<size_t N>
+xmrig::cn_hash_fun xmrig::CpuWorker<N>::fn(const Algorithm &algorithm) const {
+#if 1
+    if (!algorithm.isValid()) {
+        return nullptr;
+    }
+
+    const auto it = m_hash_fns.find(algorithm.id());
+    if (it != m_hash_fns.end()) {
+        return it->second;
+    }
+    return nullptr;
+#else
+    return CnHash::fn(algorithm, m_av, m_assembly);
+#endif
+}
+
 
 #ifdef XMRIG_ALGO_RANDOMX
 template<size_t N>
@@ -159,7 +182,7 @@ template<size_t N>
 bool xmrig::CpuWorker<N>::selfTest()
 {
 #   ifdef XMRIG_ALGO_RANDOMX
-    if (m_algorithm.family() == Algorithm::RANDOM_X) {
+    if (m_algorithm_family == Algorithm::RANDOM_X) {
         return N == 1;
     }
 #   endif
@@ -167,12 +190,12 @@ bool xmrig::CpuWorker<N>::selfTest()
     allocateCnCtx();
 
 #   ifdef XMRIG_ALGO_GHOSTRIDER
-    if (m_algorithm.family() == Algorithm::GHOSTRIDER) {
+    if (m_algorithm_family == Algorithm::GHOSTRIDER) {
         return (N == 8) && verify(Algorithm::GHOSTRIDER_RTM, test_output_gr);
     }
 #   endif
 
-    if (m_algorithm.family() == Algorithm::CN) {
+    if (m_algorithm_family == Algorithm::CN) {
         const bool rc = verify(Algorithm::CN_0,      test_output_v0)   &&
                         verify(Algorithm::CN_1,      test_output_v1)   &&
                         verify(Algorithm::CN_2,      test_output_v2)   &&
@@ -190,14 +213,14 @@ bool xmrig::CpuWorker<N>::selfTest()
     }
 
 #   ifdef XMRIG_ALGO_CN_LITE
-    if (m_algorithm.family() == Algorithm::CN_LITE) {
+    if (m_algorithm_family == Algorithm::CN_LITE) {
         return verify(Algorithm::CN_LITE_0,    test_output_v0_lite) &&
                verify(Algorithm::CN_LITE_1,    test_output_v1_lite);
     }
 #   endif
 
 #   ifdef XMRIG_ALGO_CN_HEAVY
-    if (m_algorithm.family() == Algorithm::CN_HEAVY) {
+    if (m_algorithm_family == Algorithm::CN_HEAVY) {
         return verify(Algorithm::CN_HEAVY_0,    test_output_v0_heavy)  &&
                verify(Algorithm::CN_HEAVY_XHV,  test_output_xhv_heavy) &&
                verify(Algorithm::CN_HEAVY_TUBE, test_output_tube_heavy);
@@ -205,20 +228,20 @@ bool xmrig::CpuWorker<N>::selfTest()
 #   endif
 
 #   ifdef XMRIG_ALGO_CN_PICO
-    if (m_algorithm.family() == Algorithm::CN_PICO) {
+    if (m_algorithm_family == Algorithm::CN_PICO) {
         return verify(Algorithm::CN_PICO_0, test_output_pico_trtl) &&
                verify(Algorithm::CN_PICO_TLO, test_output_pico_tlo);
     }
 #   endif
 
 #   ifdef XMRIG_ALGO_CN_FEMTO
-    if (m_algorithm.family() == Algorithm::CN_FEMTO) {
+    if (m_algorithm_family == Algorithm::CN_FEMTO) {
         return verify(Algorithm::CN_UPX2, test_output_femto_upx2);
     }
 #   endif
 
 #   ifdef XMRIG_ALGO_ARGON2
-    if (m_algorithm.family() == Algorithm::ARGON2) {
+    if (m_algorithm_family == Algorithm::ARGON2) {
         return verify(Algorithm::AR2_CHUKWA, argon2_chukwa_test_out) &&
                verify(Algorithm::AR2_CHUKWA_V2, argon2_chukwa_v2_test_out) &&
                verify(Algorithm::AR2_WRKZ, argon2_wrkz_test_out);
@@ -262,7 +285,7 @@ void xmrig::CpuWorker<N>::start()
         while (!Nonce::isOutdated(Nonce::CPU, m_job.sequence())) {
             const Job &job = m_job.currentJob();
 
-            if (job.algorithm().l3() != m_algorithm.l3()) {
+            if (job.algorithm().l3() != m_algorithm_l3) {
                 break;
             }
 
@@ -359,7 +382,9 @@ void xmrig::CpuWorker<N>::start()
             }
         }
 
-        consumeJob();
+        if (!Nonce::isPaused()) {
+            consumeJob();
+        }
     }
 }
 
@@ -487,11 +512,11 @@ void xmrig::CpuWorker<N>::allocateCnCtx()
 #       ifdef XMRIG_ALGO_CN_HEAVY
         // cn-heavy optimization for Zen3 CPUs
         if (m_memory == cn_heavyZen3Memory) {
-            shift = (id() / 8) * m_algorithm.l3() * 8 + (id() % 8) * 64;
+            shift = (id() / 8) * m_algorithm_l3 * 8 + (id() % 8) * 64;
         }
 #       endif
 
-        CnCtx::create(m_ctx, m_memory->scratchpad() + shift, m_algorithm.l3(), N);
+        CnCtx::create(m_ctx, m_memory->scratchpad() + shift, m_algorithm_l3, N);
     }
 }
 

src/backend/cpu/CpuWorker.h

@@ -26,6 +26,8 @@
 #include "base/tools/Object.h"
 #include "net/JobResult.h"
 
+#include <unordered_map>
+
 
 #ifdef XMRIG_ALGO_RANDOMX
 class randomx_vm;
@@ -55,7 +57,7 @@ public:
     size_t threads() const override
     {
 #       ifdef XMRIG_ALGO_GHOSTRIDER
-        return ((m_algorithm.family() == Algorithm::GHOSTRIDER) && m_ghHelper) ? 2 : 1;
+        return ((m_algorithm_family == Algorithm::GHOSTRIDER) && m_ghHelper) ? 2 : 1;
 #       else
         return 1;
 #       endif
@@ -71,7 +73,7 @@ protected:
     inline void jobEarlyNotification(const Job&) override   {}
 
 private:
-    inline cn_hash_fun fn(const Algorithm &algorithm) const { return CnHash::fn(algorithm, m_av, m_assembly); }
+    cn_hash_fun fn(const Algorithm &algorithm) const;
 
 #   ifdef XMRIG_ALGO_RANDOMX
     void allocateRandomX_VM();
@@ -84,8 +86,10 @@ private:
     void consumeJob();
 
     alignas(8) uint8_t m_hash[N * 32]{ 0 };
-    const Algorithm m_algorithm;
+    const uint32_t m_algorithm_family;
+    const size_t m_algorithm_l3;
     const Assembly m_assembly;
+    std::unordered_map<Algorithm::Id, cn_hash_fun> m_hash_fns;
     const bool m_hwAES;
     const bool m_yield;
     const CnHash::AlgoVariant m_av;

src/backend/cpu/interfaces/ICpuInfo.h

@@ -52,7 +52,8 @@ public:
         ARCH_ZEN_PLUS,
         ARCH_ZEN2,
         ARCH_ZEN3,
-        ARCH_ZEN4
+        ARCH_ZEN4,
+        ARCH_ZEN5
     };
 
     enum MsrMod : uint32_t {
@@ -60,12 +61,13 @@ public:
         MSR_MOD_RYZEN_17H,
         MSR_MOD_RYZEN_19H,
         MSR_MOD_RYZEN_19H_ZEN4,
+        MSR_MOD_RYZEN_1AH_ZEN5,
         MSR_MOD_INTEL,
         MSR_MOD_CUSTOM,
         MSR_MOD_MAX
     };
 
-#   define MSR_NAMES_LIST "none", "ryzen_17h", "ryzen_19h", "ryzen_19h_zen4", "intel", "custom"
+#   define MSR_NAMES_LIST "none", "ryzen_17h", "ryzen_19h", "ryzen_19h_zen4", "ryzen_1Ah_zen5", "intel", "custom"
 
     enum Flag : uint32_t {
         FLAG_AES,

src/backend/cpu/platform/BasicCpuInfo.cpp

@@ -64,7 +64,7 @@ static_assert(kCpuFlagsSize == ICpuInfo::FLAG_MAX, "kCpuFlagsSize and FLAG_MAX m
 
 
 #ifdef XMRIG_FEATURE_MSR
-constexpr size_t kMsrArraySize                                  = 6;
+constexpr size_t kMsrArraySize                                  = 7;
 static const std::array<const char *, kMsrArraySize> msrNames   = { MSR_NAMES_LIST };
 static_assert(kMsrArraySize == ICpuInfo::MSR_MOD_MAX, "kMsrArraySize and MSR_MOD_MAX mismatch");
 #endif
@@ -260,6 +260,11 @@ xmrig::BasicCpuInfo::BasicCpuInfo() :
                     }
                     break;
 
+                case 0x1a:
+                    m_arch = ARCH_ZEN5;
+                    m_msrMod = MSR_MOD_RYZEN_1AH_ZEN5;
+                    break;
+
                 default:
                     m_msrMod = MSR_MOD_NONE;
                     break;

src/backend/cpu/platform/HwlocCpuInfo.cpp

@@ -326,7 +326,8 @@ void xmrig::HwlocCpuInfo::processTopLevelCache(hwloc_obj_t cache, const Algorith
         }
     }
 
-    if (scratchpad == 2 * oneMiB) {
+    // This code is supposed to run only on Intel CPUs
+    if ((vendor() == VENDOR_INTEL) && (scratchpad == 2 * oneMiB)) {
         if (L2 && (cores.size() * oneMiB) == L2 && L2_associativity == 16 && L3 >= L2) {
             L3    = L2;
             extra = L2;

src/backend/cuda/CudaConfig_gen.h

@@ -114,7 +114,6 @@ size_t inline generate<Algorithm::RANDOM_X>(Threads<CudaThreads> &threads, const
     auto rx  = CudaThreads(devices, Algorithm::RX_0);
     auto wow = CudaThreads(devices, Algorithm::RX_WOW);
     auto arq = CudaThreads(devices, Algorithm::RX_ARQ);
-    auto kva = CudaThreads(devices, Algorithm::RX_KEVA);
 
     if (!threads.isExist(Algorithm::RX_WOW) && wow != rx) {
         count += threads.move(Algorithm::kRX_WOW, std::move(wow));
@@ -124,10 +123,6 @@ size_t inline generate<Algorithm::RANDOM_X>(Threads<CudaThreads> &threads, const
         count += threads.move(Algorithm::kRX_ARQ, std::move(arq));
     }
 
-    if (!threads.isExist(Algorithm::RX_KEVA) && kva != rx) {
-        count += threads.move(Algorithm::kRX_KEVA, std::move(kva));
-    }
-
     count += threads.move(Algorithm::kRX, std::move(rx));
 
     return count;

src/backend/cuda/CudaWorker.cpp

@@ -158,7 +158,7 @@ void xmrig::CudaWorker::start()
             std::this_thread::yield();
         }
 
-        if (!consumeJob()) {
+        if (isReady() && !consumeJob()) {
             return;
         }
     }

src/backend/opencl/OclWorker.cpp

@@ -170,7 +170,7 @@ void xmrig::OclWorker::start()
             const uint64_t t = Chrono::steadyMSecs();
 
             try {
-                m_runner->run(readUnaligned(m_job.nonce()), results);
+                m_runner->run(readUnaligned(m_job.nonce()), m_job.nonceOffset(), results);
             }
             catch (std::exception &ex) {
                 printError(id(), ex.what());
@@ -190,7 +190,7 @@ void xmrig::OclWorker::start()
             std::this_thread::yield();
         }
 
-        if (!consumeJob()) {
+        if (isReady() && !consumeJob()) {
             return;
         }
     }

src/backend/opencl/cl/cn/algorithm.cl

@@ -22,8 +22,8 @@
 #define ALGO_RX_WOW         0x72141177
 #define ALGO_RX_ARQMA       0x72121061
 #define ALGO_RX_SFX         0x72151273
-#define ALGO_RX_KEVA        0x7214116b
 #define ALGO_RX_GRAFT       0x72151267
+#define ALGO_RX_YADA        0x72151279
 #define ALGO_AR2_CHUKWA     0x61130000
 #define ALGO_AR2_CHUKWA_V2  0x61140000
 #define ALGO_AR2_WRKZ       0x61120000

src/backend/opencl/cl/cn/cryptonight_cl.h

@@ -34,9 +34,9 @@ static const char cryptonight_cl[61447] = {
     0x31,0x35,0x31,0x32,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x57,0x4f,0x57,0x20,0x30,0x78,0x37,0x32,0x31,
     0x34,0x31,0x31,0x37,0x37,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x41,0x52,0x51,0x4d,0x41,0x20,0x30,0x78,0x37,0x32,
     0x31,0x32,0x31,0x30,0x36,0x31,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x53,0x46,0x58,0x20,0x30,0x78,0x37,0x32,0x31,
-    0x35,0x31,0x32,0x37,0x33,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x4b,0x45,0x56,0x41,0x20,0x30,0x78,0x37,0x32,0x31,
-    0x34,0x31,0x31,0x36,0x62,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x47,0x52,0x41,0x46,0x54,0x20,0x30,0x78,0x37,0x32,
-    0x31,0x35,0x31,0x32,0x36,0x37,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,0x41,0x20,0x30,
+    0x35,0x31,0x32,0x37,0x33,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x47,0x52,0x41,0x46,0x54,0x20,0x30,0x78,0x37,0x32,
+    0x31,0x35,0x31,0x32,0x36,0x37,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x59,0x41,0x44,0x41,0x20,0x30,0x78,0x37,0x32,
+    0x31,0x35,0x31,0x32,0x37,0x39,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,0x41,0x20,0x30,
     0x78,0x36,0x31,0x31,0x33,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,
     0x41,0x5f,0x56,0x32,0x20,0x30,0x78,0x36,0x31,0x31,0x34,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,
     0x5f,0x57,0x52,0x4b,0x5a,0x20,0x30,0x78,0x36,0x31,0x31,0x32,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x4b,0x41,

src/backend/opencl/cl/rx/blake2b.cl

@@ -225,6 +225,110 @@ __kernel void blake2b_initial_hash_double(__global void *out, __global const voi
 	t[7] = hash[7];
 }
 
+void blake2b_512_process_big_block(ulong *out, __global const ulong* in, uint in_len, uint out_len, uint nonce, uint nonce_offset)
+{
+	ulong h[8] = { iv0 ^ (0x01010000u | out_len), iv1, iv2, iv3, iv4, iv5, iv6, iv7 };
+
+	for (uint t = 128; t < in_len; t += 128, in += 16) {
+		ulong m[16] = { in[0], in[1], in[2], in[3], in[4], in[5], in[6], in[7], in[8], in[9], in[10], in[11],  in[12], in[13], in[14], in[15] };
+
+		const uint k0 = (nonce_offset + 0) - (t - 128);
+		const uint k1 = (nonce_offset + 1) - (t - 128);
+		const uint k2 = (nonce_offset + 2) - (t - 128);
+		const uint k3 = (nonce_offset + 3) - (t - 128);
+
+		if (k0 < 128) m[k0 / 8] |= (ulong)((nonce >>  0) & 255) << ((k0 % 8) * 8);
+		if (k1 < 128) m[k1 / 8] |= (ulong)((nonce >>  8) & 255) << ((k1 % 8) * 8);
+		if (k2 < 128) m[k2 / 8] |= (ulong)((nonce >> 16) & 255) << ((k2 % 8) * 8);
+		if (k3 < 128) m[k3 / 8] |= (ulong)((nonce >> 24) & 255) << ((k3 % 8) * 8);
+
+		ulong v[16] = {  h[0],  h[1],  h[2],  h[3],  h[4],  h[5],  h[6],  h[7],   iv0,   iv1,    iv2,    iv3, iv4 ^ t,    iv5,    iv6,    iv7 };
+
+		BLAKE2B_ROUNDS();
+
+		h[0] ^= v[0] ^ v[ 8];
+		h[1] ^= v[1] ^ v[ 9];
+		h[2] ^= v[2] ^ v[10];
+		h[3] ^= v[3] ^ v[11];
+		h[4] ^= v[4] ^ v[12];
+		h[5] ^= v[5] ^ v[13];
+		h[6] ^= v[6] ^ v[14];
+		h[7] ^= v[7] ^ v[15];
+	}
+
+	uint k = in_len & 127;
+	if (k == 0) k = 128;
+
+	ulong m[16] = {
+		(k >   0) ? in[ 0] : 0,
+		(k >   8) ? in[ 1] : 0,
+		(k >  16) ? in[ 2] : 0,
+		(k >  24) ? in[ 3] : 0,
+		(k >  32) ? in[ 4] : 0,
+		(k >  40) ? in[ 5] : 0,
+		(k >  48) ? in[ 6] : 0,
+		(k >  56) ? in[ 7] : 0,
+		(k >  64) ? in[ 8] : 0,
+		(k >  72) ? in[ 9] : 0,
+		(k >  80) ? in[10] : 0,
+		(k >  88) ? in[11] : 0,
+		(k >  96) ? in[12] : 0,
+		(k > 104) ? in[13] : 0,
+		(k > 112) ? in[14] : 0,
+		(k > 120) ? in[15] : 0
+	};
+
+	const uint t = in_len - k;
+
+	const uint k0 = nonce_offset + 0 - t;
+	const uint k1 = nonce_offset + 1 - t;
+	const uint k2 = nonce_offset + 2 - t;
+	const uint k3 = nonce_offset + 3 - t;
+
+	if (k0 < k) m[k0 / 8] |= (ulong)((nonce >>  0) & 255) << ((k0 % 8) * 8);
+	if (k1 < k) m[k1 / 8] |= (ulong)((nonce >>  8) & 255) << ((k1 % 8) * 8);
+	if (k2 < k) m[k2 / 8] |= (ulong)((nonce >> 16) & 255) << ((k2 % 8) * 8);
+	if (k3 < k) m[k3 / 8] |= (ulong)((nonce >> 24) & 255) << ((k3 % 8) * 8);
+
+	if (k % 8) {
+		m[k / 8] &= (ulong)(-1) >> (64 - (k % 8) * 8);
+	}
+
+	ulong v[16] = { h[0], h[1], h[2], h[3], h[4], h[5], h[6], h[7], iv0, iv1, iv2, iv3, iv4 ^ in_len, iv5, ~iv6, iv7 };
+
+	BLAKE2B_ROUNDS();
+
+	if (out_len >  0) out[0] = h[0] ^ v[0] ^ v[8];
+	if (out_len >  8) out[1] = h[1] ^ v[1] ^ v[9];
+	if (out_len > 16) out[2] = h[2] ^ v[2] ^ v[10];
+	if (out_len > 24) out[3] = h[3] ^ v[3] ^ v[11];
+	if (out_len > 32) out[4] = h[4] ^ v[4] ^ v[12];
+	if (out_len > 40) out[5] = h[5] ^ v[5] ^ v[13];
+	if (out_len > 48) out[6] = h[6] ^ v[6] ^ v[14];
+	if (out_len > 56) out[7] = h[7] ^ v[7] ^ v[15];
+}
+
+__attribute__((reqd_work_group_size(64, 1, 1)))
+__kernel void blake2b_initial_hash_big(__global void *out, __global const void* blockTemplate, uint blockTemplateSize, uint start_nonce, uint nonce_offset)
+{
+	const uint global_index = get_global_id(0);
+
+	__global const ulong* p = (__global const ulong*) blockTemplate;
+
+	ulong hash[8];
+	blake2b_512_process_big_block(hash, p, blockTemplateSize, 64, start_nonce + global_index, nonce_offset);
+
+	__global ulong* t = ((__global ulong*) out) + global_index * 8;
+	t[0] = hash[0];
+	t[1] = hash[1];
+	t[2] = hash[2];
+	t[3] = hash[3];
+	t[4] = hash[4];
+	t[5] = hash[5];
+	t[6] = hash[6];
+	t[7] = hash[7];
+}
+
 #define in_len 256
 
 #define out_len 32

src/backend/opencl/cl/rx/randomx.cl

@@ -1,13 +1,11 @@
 #include "../cn/algorithm.cl"
 
-#if (ALGO == ALGO_RX_0)
+#if ((ALGO == ALGO_RX_0) || (ALGO == ALGO_RX_YADA))
 #include "randomx_constants_monero.h"
 #elif (ALGO == ALGO_RX_WOW)
 #include "randomx_constants_wow.h"
 #elif (ALGO == ALGO_RX_ARQMA)
 #include "randomx_constants_arqma.h"
-#elif (ALGO == ALGO_RX_KEVA)
-#include "randomx_constants_keva.h"
 #elif (ALGO == ALGO_RX_GRAFT)
 #include "randomx_constants_graft.h"
 #endif

src/backend/opencl/cl/rx/randomx_constants_keva.h

@@ -1,96 +0,0 @@
-/*
-Copyright (c) 2019 SChernykh
-
-This file is part of RandomX OpenCL.
-
-RandomX OpenCL is free software: you can redistribute it and/or modify
-it under the terms of the GNU General Public License as published by
-the Free Software Foundation, either version 3 of the License, or
-(at your option) any later version.
-
-RandomX OpenCL is distributed in the hope that it will be useful,
-but WITHOUT ANY WARRANTY; without even the implied warranty of
-MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
-GNU General Public License for more details.
-
-You should have received a copy of the GNU General Public License
-along with RandomX OpenCL. If not, see <http://www.gnu.org/licenses/>.
-*/
-
-//Dataset base size in bytes. Must be a power of 2.
-#define RANDOMX_DATASET_BASE_SIZE  2147483648
-
-//Dataset extra size. Must be divisible by 64.
-#define RANDOMX_DATASET_EXTRA_SIZE 33554368
-
-//Scratchpad L3 size in bytes. Must be a power of 2.
-#define RANDOMX_SCRATCHPAD_L3      1048576
-
-//Scratchpad L2 size in bytes. Must be a power of two and less than or equal to RANDOMX_SCRATCHPAD_L3.
-#define RANDOMX_SCRATCHPAD_L2      131072
-
-//Scratchpad L1 size in bytes. Must be a power of two (minimum 64) and less than or equal to RANDOMX_SCRATCHPAD_L2.
-#define RANDOMX_SCRATCHPAD_L1      16384
-
-//Jump condition mask size in bits.
-#define RANDOMX_JUMP_BITS          8
-
-//Jump condition mask offset in bits. The sum of RANDOMX_JUMP_BITS and RANDOMX_JUMP_OFFSET must not exceed 16.
-#define RANDOMX_JUMP_OFFSET        8
-
-//Integer instructions
-#define RANDOMX_FREQ_IADD_RS       16
-#define RANDOMX_FREQ_IADD_M         7
-#define RANDOMX_FREQ_ISUB_R        16
-#define RANDOMX_FREQ_ISUB_M         7
-#define RANDOMX_FREQ_IMUL_R        16
-#define RANDOMX_FREQ_IMUL_M         4
-#define RANDOMX_FREQ_IMULH_R        4
-#define RANDOMX_FREQ_IMULH_M        1
-#define RANDOMX_FREQ_ISMULH_R       4
-#define RANDOMX_FREQ_ISMULH_M       1
-#define RANDOMX_FREQ_IMUL_RCP       8
-#define RANDOMX_FREQ_INEG_R         2
-#define RANDOMX_FREQ_IXOR_R        15
-#define RANDOMX_FREQ_IXOR_M         5
-#define RANDOMX_FREQ_IROR_R         8
-#define RANDOMX_FREQ_IROL_R         2
-#define RANDOMX_FREQ_ISWAP_R        4
-
-//Floating point instructions
-#define RANDOMX_FREQ_FSWAP_R        4
-#define RANDOMX_FREQ_FADD_R        16
-#define RANDOMX_FREQ_FADD_M         5
-#define RANDOMX_FREQ_FSUB_R        16
-#define RANDOMX_FREQ_FSUB_M         5
-#define RANDOMX_FREQ_FSCAL_R        6
-#define RANDOMX_FREQ_FMUL_R        32
-#define RANDOMX_FREQ_FDIV_M         4
-#define RANDOMX_FREQ_FSQRT_R        6
-
-//Control instructions
-#define RANDOMX_FREQ_CBRANCH       25
-#define RANDOMX_FREQ_CFROUND        1
-
-//Store instruction
-#define RANDOMX_FREQ_ISTORE        16
-
-//No-op instruction
-#define RANDOMX_FREQ_NOP            0
-
-#define RANDOMX_DATASET_ITEM_SIZE 64
-
-#define RANDOMX_PROGRAM_SIZE 256
-
-#define HASH_SIZE 64
-#define ENTROPY_SIZE (128 + RANDOMX_PROGRAM_SIZE * 8)
-#define REGISTERS_SIZE 256
-#define IMM_BUF_SIZE (RANDOMX_PROGRAM_SIZE * 4 - REGISTERS_SIZE)
-#define IMM_INDEX_COUNT ((IMM_BUF_SIZE / 4) - 2)
-#define VM_STATE_SIZE (REGISTERS_SIZE + IMM_BUF_SIZE + RANDOMX_PROGRAM_SIZE * 4)
-#define ROUNDING_MODE (RANDOMX_FREQ_CFROUND ? -1 : 0)
-
-// Scratchpad L1/L2/L3 bits
-#define LOC_L1 (32 - 14)
-#define LOC_L2 (32 - 17)
-#define LOC_L3 (32 - 20)

src/backend/opencl/interfaces/IOclRunner.h

@@ -64,7 +64,7 @@ public:
     virtual uint32_t deviceIndex() const                    = 0;
     virtual void build()                                    = 0;
     virtual void init()                                     = 0;
-    virtual void run(uint32_t nonce, uint32_t *hashOutput)  = 0;
+    virtual void run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput)  = 0;
     virtual void set(const Job &job, uint8_t *blob)         = 0;
     virtual void jobEarlyNotification(const Job&)           = 0;
 

src/backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.cpp

@@ -0,0 +1,59 @@
+/* XMRig
+ * Copyright 2010      Jeff Garzik <jgarzik@pobox.com>
+ * Copyright 2012-2014 pooler      <pooler@litecoinpool.org>
+ * Copyright 2014      Lucas Jones <https://github.com/lucasjones>
+ * Copyright 2014-2016 Wolf9466    <https://github.com/OhGodAPet>
+ * Copyright 2016      Jay D Dee   <jayddee246@gmail.com>
+ * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
+ * Copyright 2018-2019 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2019 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ *
+ *   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
+ *   the Free Software Foundation, either version 3 of the License, or
+ *   (at your option) any later version.
+ *
+ *   This program is distributed in the hope that it will be useful,
+ *   but WITHOUT ANY WARRANTY; without even the implied warranty of
+ *   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ *   GNU General Public License for more details.
+ *
+ *   You should have received a copy of the GNU General Public License
+ *   along with this program. If not, see <http://www.gnu.org/licenses/>.
+ */
+
+
+#include "backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+void xmrig::Blake2bInitialHashBigKernel::enqueue(cl_command_queue queue, size_t threads)
+{
+    const size_t gthreads        = threads;
+    static const size_t lthreads = 64;
+
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+}
+
+
+// __kernel void blake2b_initial_hash_double(__global void *out, __global const void* blockTemplate, uint blockTemplateSize, uint start_nonce)
+void xmrig::Blake2bInitialHashBigKernel::setArgs(cl_mem out, cl_mem blockTemplate)
+{
+    setArg(0, sizeof(cl_mem), &out);
+    setArg(1, sizeof(cl_mem), &blockTemplate);
+}
+
+
+void xmrig::Blake2bInitialHashBigKernel::setBlobSize(size_t size)
+{
+    const uint32_t s = size;
+
+    setArg(2, sizeof(uint32_t), &s);
+}
+
+
+void xmrig::Blake2bInitialHashBigKernel::setNonce(uint32_t nonce, uint32_t nonce_offset)
+{
+    setArg(3, sizeof(uint32_t), &nonce);
+    setArg(4, sizeof(uint32_t), &nonce_offset);
+}

src/backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.h

@@ -0,0 +1,50 @@
+/* XMRig
+ * Copyright 2010      Jeff Garzik <jgarzik@pobox.com>
+ * Copyright 2012-2014 pooler      <pooler@litecoinpool.org>
+ * Copyright 2014      Lucas Jones <https://github.com/lucasjones>
+ * Copyright 2014-2016 Wolf9466    <https://github.com/OhGodAPet>
+ * Copyright 2016      Jay D Dee   <jayddee246@gmail.com>
+ * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
+ * Copyright 2018-2019 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2019 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ *
+ *   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
+ *   the Free Software Foundation, either version 3 of the License, or
+ *   (at your option) any later version.
+ *
+ *   This program is distributed in the hope that it will be useful,
+ *   but WITHOUT ANY WARRANTY; without even the implied warranty of
+ *   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ *   GNU General Public License for more details.
+ *
+ *   You should have received a copy of the GNU General Public License
+ *   along with this program. If not, see <http://www.gnu.org/licenses/>.
+ */
+
+#ifndef XMRIG_BLAKE2BINITIALHASHBIGKERNEL_H
+#define XMRIG_BLAKE2BINITIALHASHBIGKERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+class Blake2bInitialHashBigKernel : public OclKernel
+{
+public:
+    inline Blake2bInitialHashBigKernel(cl_program program) : OclKernel(program, "blake2b_initial_hash_big") {}
+
+    void enqueue(cl_command_queue queue, size_t threads);
+    void setArgs(cl_mem out, cl_mem blockTemplate);
+    void setBlobSize(size_t size);
+    void setNonce(uint32_t nonce, uint32_t nonce_offset);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_BLAKE2BINITIALHASHBIGKERNEL_H */

src/backend/opencl/opencl.cmake

@@ -80,6 +80,7 @@ if (WITH_OPENCL)
     if (WITH_RANDOMX)
         list(APPEND HEADERS_BACKEND_OPENCL
              src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h
+             src/backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.h
              src/backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.h
              src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.h
              src/backend/opencl/kernels/rx/ExecuteVmKernel.h
@@ -97,6 +98,7 @@ if (WITH_OPENCL)
         list(APPEND SOURCES_BACKEND_OPENCL
              src/backend/opencl/generators/ocl_generic_rx_generator.cpp
              src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.cpp
+             src/backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.cpp
              src/backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.cpp
              src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.cpp
              src/backend/opencl/kernels/rx/ExecuteVmKernel.cpp

src/backend/opencl/runners/OclCnRunner.cpp

@@ -87,7 +87,7 @@ size_t xmrig::OclCnRunner::bufferSize() const
 }
 
 
-void xmrig::OclCnRunner::run(uint32_t nonce, uint32_t *hashOutput)
+void xmrig::OclCnRunner::run(uint32_t nonce, uint32_t /*nonce_offset*/, uint32_t *hashOutput)
 {
     static const cl_uint zero = 0;
 

src/backend/opencl/runners/OclCnRunner.h

@@ -42,7 +42,7 @@ public:
 
 protected:
     size_t bufferSize() const override;
-    void run(uint32_t nonce, uint32_t *hashOutput) override;
+    void run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput) override;
     void set(const Job &job, uint8_t *blob) override;
     void build() override;
     void init() override;

src/backend/opencl/runners/OclKawPowRunner.cpp

@@ -75,7 +75,7 @@ OclKawPowRunner::~OclKawPowRunner()
 }
 
 
-void OclKawPowRunner::run(uint32_t nonce, uint32_t *hashOutput)
+void OclKawPowRunner::run(uint32_t nonce, uint32_t /*nonce_offset*/, uint32_t *hashOutput)
 {
     const size_t local_work_size = m_workGroupSize;
     const size_t global_work_offset = nonce;

src/backend/opencl/runners/OclKawPowRunner.h

@@ -40,7 +40,7 @@ public:
     ~OclKawPowRunner() override;
 
 protected:
-    void run(uint32_t nonce, uint32_t *hashOutput) override;
+    void run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput) override;
     void set(const Job &job, uint8_t *blob) override;
     void build() override;
     void init() override;

src/backend/opencl/runners/OclRxBaseRunner.cpp

@@ -26,6 +26,7 @@
 #include "backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h"
 #include "backend/opencl/kernels/rx/Blake2bInitialHashKernel.h"
 #include "backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.h"
+#include "backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.h"
 #include "backend/opencl/kernels/rx/FillAesKernel.h"
 #include "backend/opencl/kernels/rx/FindSharesKernel.h"
 #include "backend/opencl/kernels/rx/HashAesKernel.h"
@@ -73,6 +74,7 @@ xmrig::OclRxBaseRunner::~OclRxBaseRunner()
     delete m_hashAes1Rx4;
     delete m_blake2b_initial_hash;
     delete m_blake2b_initial_hash_double;
+    delete m_blake2b_initial_hash_big;
     delete m_blake2b_hash_registers_32;
     delete m_blake2b_hash_registers_64;
     delete m_find_shares;
@@ -85,7 +87,7 @@ xmrig::OclRxBaseRunner::~OclRxBaseRunner()
 }
 
 
-void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
+void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput)
 {
     static const uint32_t zero = 0;
 
@@ -96,8 +98,7 @@ void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
         m_blake2b_initial_hash_double->setNonce(nonce);
     }
     else {
-        hashOutput[0xFF] = 0;
-        return;
+        m_blake2b_initial_hash_big->setNonce(nonce, nonce_offset);
     }
 
     m_find_shares->setNonce(nonce);
@@ -107,9 +108,12 @@ void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
     if (m_jobSize <= 128) {
         m_blake2b_initial_hash->enqueue(m_queue, m_intensity);
     }
-    else {
+    else if (m_jobSize <= 256) {
         m_blake2b_initial_hash_double->enqueue(m_queue, m_intensity);
     }
+    else {
+        m_blake2b_initial_hash_big->enqueue(m_queue, m_intensity);
+    }
 
     m_fillAes1Rx4_scratchpad->enqueue(m_queue, m_intensity);
 
@@ -150,12 +154,15 @@ void xmrig::OclRxBaseRunner::set(const Job &job, uint8_t *blob)
         memset(blob + job.size(), 0, Job::kMaxBlobSize - job.size());
     }
 
+    memset(blob + job.nonceOffset(), 0, job.nonceSize());
+
     enqueueWriteBuffer(m_input, CL_TRUE, 0, Job::kMaxBlobSize, blob);
 
     m_jobSize = job.size();
 
     m_blake2b_initial_hash->setBlobSize(job.size());
     m_blake2b_initial_hash_double->setBlobSize(job.size());
+    m_blake2b_initial_hash_big->setBlobSize(job.size());
 
     m_find_shares->setTarget(job.target());
 }
@@ -191,6 +198,9 @@ void xmrig::OclRxBaseRunner::build()
     m_blake2b_initial_hash_double = new Blake2bInitialHashDoubleKernel(m_program);
     m_blake2b_initial_hash_double->setArgs(m_hashes, m_input);
 
+    m_blake2b_initial_hash_big = new Blake2bInitialHashBigKernel(m_program);
+    m_blake2b_initial_hash_big->setArgs(m_hashes, m_input);
+
     m_blake2b_hash_registers_32 = new Blake2bHashRegistersKernel(m_program, "blake2b_hash_registers_32");
     m_blake2b_hash_registers_64 = new Blake2bHashRegistersKernel(m_program, "blake2b_hash_registers_64");
 

src/backend/opencl/runners/OclRxBaseRunner.h

@@ -36,6 +36,7 @@ namespace xmrig {
 class Blake2bHashRegistersKernel;
 class Blake2bInitialHashKernel;
 class Blake2bInitialHashDoubleKernel;
+class Blake2bInitialHashBigKernel;
 class FillAesKernel;
 class FindSharesKernel;
 class HashAesKernel;
@@ -53,7 +54,7 @@ protected:
     size_t bufferSize() const override;
     void build() override;
     void init() override;
-    void run(uint32_t nonce, uint32_t *hashOutput) override;
+    void run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput) override;
     void set(const Job &job, uint8_t *blob) override;
 
 protected:
@@ -63,6 +64,7 @@ protected:
     Blake2bHashRegistersKernel *m_blake2b_hash_registers_64       = nullptr;
     Blake2bInitialHashKernel *m_blake2b_initial_hash              = nullptr;
     Blake2bInitialHashDoubleKernel *m_blake2b_initial_hash_double = nullptr;
+    Blake2bInitialHashBigKernel* m_blake2b_initial_hash_big       = nullptr;
     Buffer m_seed;
     cl_mem m_dataset                                              = nullptr;
     cl_mem m_entropy                                              = nullptr;

src/base/api/Api.cpp

@@ -1,6 +1,6 @@
 /* XMRig
- * Copyright (c) 2018-2023 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2023 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright (c) 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright (c) 2016-2024 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
  *
  *   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
@@ -25,6 +25,8 @@
 #include "base/crypto/keccak.h"
 #include "base/io/Env.h"
 #include "base/io/json/Json.h"
+#include "base/io/log/Log.h"
+#include "base/io/log/Tags.h"
 #include "base/kernel/Base.h"
 #include "base/tools/Chrono.h"
 #include "base/tools/Cvt.h"
@@ -39,7 +41,6 @@
 
 
 #include <thread>
-#include <iostream>
 
 
 namespace xmrig {
@@ -81,8 +82,7 @@ static rapidjson::Value getResources(rapidjson::Document &doc)
 
 xmrig::Api::Api(Base *base) :
     m_base(base),
-    m_timestamp(Chrono::currentMSecsSinceEpoch()),
-    m_httpd(nullptr)
+    m_timestamp(Chrono::currentMSecsSinceEpoch())
 {
     base->addListener(this);
 
@@ -118,7 +118,8 @@ void xmrig::Api::start()
     if (!m_httpd) {
         m_httpd = new Httpd(m_base);
         if (!m_httpd->start()) {
-            std::cerr << "HTTP server failed to start." << std::endl;
+            LOG_ERR("%s " RED_BOLD("HTTP API server failed to start."), Tags::network());
+
             delete m_httpd; // Properly handle failure to start
             m_httpd = nullptr;
         }

src/base/api/Api.h

@@ -1,6 +1,6 @@
 /* XMRig
- * Copyright (c) 2018-2023 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2023 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright (c) 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright (c) 2016-2024 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
  *
  *   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

src/base/crypto/Algorithm.cpp

@@ -81,7 +81,7 @@ const char *Algorithm::kRX_WOW          = "rx/wow";
 const char *Algorithm::kRX_ARQ          = "rx/arq";
 const char *Algorithm::kRX_GRAFT        = "rx/graft";
 const char *Algorithm::kRX_SFX          = "rx/sfx";
-const char *Algorithm::kRX_KEVA         = "rx/keva";
+const char *Algorithm::kRX_YADA         = "rx/yada";
 #endif
 
 #ifdef XMRIG_ALGO_ARGON2
@@ -147,7 +147,7 @@ static const std::map<uint32_t, const char *> kAlgorithmNames = {
     ALGO_NAME(RX_ARQ),
     ALGO_NAME(RX_GRAFT),
     ALGO_NAME(RX_SFX),
-    ALGO_NAME(RX_KEVA),
+    ALGO_NAME(RX_YADA),
 #   endif
 
 #   ifdef XMRIG_ALGO_ARGON2
@@ -261,8 +261,8 @@ static const std::map<const char *, Algorithm::Id, aliasCompare> kAlgorithmAlias
                                     ALGO_ALIAS(RX_GRAFT,        "randomgraft"),
     ALGO_ALIAS_AUTO(RX_SFX),        ALGO_ALIAS(RX_SFX,          "randomx/sfx"),
                                     ALGO_ALIAS(RX_SFX,          "randomsfx"),
-    ALGO_ALIAS_AUTO(RX_KEVA),       ALGO_ALIAS(RX_KEVA,         "randomx/keva"),
-                                    ALGO_ALIAS(RX_KEVA,         "randomkeva"),
+    ALGO_ALIAS_AUTO(RX_YADA),       ALGO_ALIAS(RX_YADA,         "randomx/yada"),
+                                    ALGO_ALIAS(RX_YADA,         "randomyada"),
 #   endif
 
 #   ifdef XMRIG_ALGO_ARGON2
@@ -350,7 +350,7 @@ std::vector<xmrig::Algorithm> xmrig::Algorithm::all(const std::function<bool(con
         CN_HEAVY_0, CN_HEAVY_TUBE, CN_HEAVY_XHV,
         CN_PICO_0, CN_PICO_TLO,
         CN_UPX2,
-        RX_0, RX_WOW, RX_ARQ, RX_GRAFT, RX_SFX, RX_KEVA,
+        RX_0, RX_WOW, RX_ARQ, RX_GRAFT, RX_SFX, RX_YADA,
         AR2_CHUKWA, AR2_CHUKWA_V2, AR2_WRKZ,
         KAWPOW_RVN,
         GHOSTRIDER_RTM

src/base/crypto/Algorithm.h

@@ -77,7 +77,7 @@ public:
         RX_ARQ          = 0x72121061,   // "rx/arq"           RandomARQ (Arqma).
         RX_GRAFT        = 0x72151267,   // "rx/graft"         RandomGRAFT (Graft).
         RX_SFX          = 0x72151273,   // "rx/sfx"           RandomSFX (Safex Cash).
-        RX_KEVA         = 0x7214116b,   // "rx/keva"          RandomKEVA (Keva).
+        RX_YADA         = 0x72151279,   // "rx/yada"          RandomYada (YadaCoin).
         AR2_CHUKWA      = 0x61130000,   // "argon2/chukwa"    Argon2id (Chukwa).
         AR2_CHUKWA_V2   = 0x61140000,   // "argon2/chukwav2"  Argon2id (Chukwa v2).
         AR2_WRKZ        = 0x61120000,   // "argon2/wrkz"      Argon2id (WRKZ)
@@ -143,7 +143,7 @@ public:
     static const char *kRX_ARQ;
     static const char *kRX_GRAFT;
     static const char *kRX_SFX;
-    static const char *kRX_KEVA;
+    static const char *kRX_YADA;
 #   endif
 
 #   ifdef XMRIG_ALGO_ARGON2

src/base/crypto/Coin.cpp

@@ -50,11 +50,11 @@ static const CoinInfo coinInfo[] = {
     { Algorithm::CN_R,            "SUMO",     "Sumokoin",     240,    1000000000,     BLUE_BG_BOLD(   WHITE_BOLD_S " sumo    ") },
     { Algorithm::RX_ARQ,          "ARQ",      "ArQmA",        120,    1000000000,     BLUE_BG_BOLD(   WHITE_BOLD_S " arqma   ") },
     { Algorithm::RX_GRAFT,        "GRFT",     "Graft",        120,    10000000000,    BLUE_BG_BOLD(   WHITE_BOLD_S " graft   ") },
-    { Algorithm::RX_KEVA,         "KVA",      "Kevacoin",     0,      0,              MAGENTA_BG_BOLD(WHITE_BOLD_S " keva    ") },
     { Algorithm::KAWPOW_RVN,      "RVN",      "Ravencoin",    0,      0,              BLUE_BG_BOLD(   WHITE_BOLD_S " raven   ") },
     { Algorithm::RX_WOW,          "WOW",      "Wownero",      300,    100000000000,   MAGENTA_BG_BOLD(WHITE_BOLD_S " wownero ") },
     { Algorithm::RX_0,            "ZEPH",     "Zephyr",       120,    1000000000000,  BLUE_BG_BOLD(   WHITE_BOLD_S " zephyr  ") },
     { Algorithm::RX_0,            "Townforge","Townforge",    30,     100000000,      MAGENTA_BG_BOLD(WHITE_BOLD_S " townforge ") },
+    { Algorithm::RX_YADA,         "YDA",      "YadaCoin",     120,    100000000,      BLUE_BG_BOLD(   WHITE_BOLD_S " yada    ") },
 };
 
 

src/base/crypto/Coin.h

@@ -36,11 +36,11 @@ public:
         SUMO,
         ARQMA,
         GRAFT,
-        KEVA,
         RAVEN,
         WOWNERO,
         ZEPHYR,
         TOWNFORGE,
+        YADA,
         MAX
     };
 

src/base/io/json/Json.cpp

@@ -211,11 +211,13 @@ rapidjson::Value xmrig::Json::normalize(double value, bool zero)
 {
     using namespace rapidjson;
 
-    if (!std::isnormal(value)) {
+    const double value_rounded = floor(value * 100.0) / 100.0;
+
+    if (!std::isnormal(value) || !std::isnormal(value_rounded)) {
         return zero ? Value(0.0) : Value(kNullType);
     }
 
-    return Value(floor(value * 100.0) / 100.0);
+    return Value(value_rounded);
 }
 
 

src/base/net/stratum/Client.cpp

@@ -363,7 +363,7 @@ bool xmrig::Client::parseJob(const rapidjson::Value &params, int *code)
 
     Job job(has<EXT_NICEHASH>(), m_pool.algorithm(), m_rpcId);
 
-    if (!job.setId(params["job_id"].GetString())) {
+    if (!job.setId(Json::getString(params, "job_id"))) {
         *code = 3;
         return false;
     }
@@ -400,7 +400,7 @@ bool xmrig::Client::parseJob(const rapidjson::Value &params, int *code)
         }
     }
 
-    if (!job.setTarget(params["target"].GetString())) {
+    if (!job.setTarget(Json::getString(params, "target"))) {
         *code = 5;
         return false;
     }

src/base/net/stratum/Job.cpp

@@ -7,8 +7,8 @@
  * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
  * Copyright 2018      Lee Clagett <https://github.com/vtnerd>
  * Copyright 2019      Howard Chu  <https://github.com/hyc>
- * Copyright 2018-2021 SChernykh   <https://github.com/SChernykh>
- * Copyright 2016-2021 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2024 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
  *
  *   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
@@ -24,11 +24,9 @@
  *   along with this program. If not, see <http://www.gnu.org/licenses/>.
  */
 
-
 #include <cassert>
 #include <cstring>
 
-
 #include "base/net/stratum/Job.h"
 #include "base/tools/Alignment.h"
 #include "base/tools/Buffer.h"
@@ -112,35 +110,69 @@ bool xmrig::Job::setSeedHash(const char *hash)
 
 bool xmrig::Job::setTarget(const char *target)
 {
-    if (!target) {
+    static auto parse = [](const char *target, size_t size, const Algorithm &algorithm) -> uint64_t {
+        if (algorithm == Algorithm::RX_YADA) {
+            return strtoull(target, nullptr, 16);
+        }
+
+        const auto raw = Cvt::fromHex(target, size);
+
+        switch (raw.size()) {
+        case 4:
+            return 0xFFFFFFFFFFFFFFFFULL / (0xFFFFFFFFULL / uint64_t(*reinterpret_cast<const uint32_t *>(raw.data())));
+
+        case 8:
+            return *reinterpret_cast<const uint64_t *>(raw.data());
+
+        default:
+            break;
+        }
+
+        return 0;
+    };
+
+    const size_t size = target ? strlen(target) : 0;
+
+    if (size < 4 || (m_target = parse(target, size, algorithm())) == 0) {
         return false;
     }
 
-    const auto raw    = Cvt::fromHex(target, strlen(target));
-    const size_t size = raw.size();
-
-    if (size == 4) {
-        m_target = 0xFFFFFFFFFFFFFFFFULL / (0xFFFFFFFFULL / uint64_t(*reinterpret_cast<const uint32_t *>(raw.data())));
-    }
-    else if (size == 8) {
-        m_target = *reinterpret_cast<const uint64_t *>(raw.data());
-    }
-    else {
-        return false;
-    }
-
-#   ifdef XMRIG_PROXY_PROJECT
-    assert(sizeof(m_rawTarget) > (size * 2));
-
-    memset(m_rawTarget, 0, sizeof(m_rawTarget));
-    memcpy(m_rawTarget, target, std::min(size * 2, sizeof(m_rawTarget)));
-#   endif
-
     m_diff = toDiff(m_target);
+
+#   ifdef XMRIG_PROXY_PROJECT
+    if (size >= sizeof(m_rawTarget)) {
+        return false;
+    }
+
+    memset(m_rawTarget, 0, sizeof(m_rawTarget));
+    memcpy(m_rawTarget, target, size);
+#   endif
+
     return true;
 }
 
 
+size_t xmrig::Job::nonceOffset() const
+{
+    switch (algorithm().family()) {
+    case Algorithm::KAWPOW:
+        return 32;
+
+    case Algorithm::GHOSTRIDER:
+        return 76;
+
+    default:
+        break;
+    }
+
+    if (algorithm() == Algorithm::RX_YADA) {
+        return 147;
+    }
+
+    return 39;
+}
+
+
 void xmrig::Job::setDiff(uint64_t diff)
 {
     m_diff   = diff;
@@ -171,14 +203,6 @@ void xmrig::Job::setSigKey(const char *sig_key)
 }
 
 
-int32_t xmrig::Job::nonceOffset() const
-{
-   auto f = algorithm().family();
-   if (f == Algorithm::KAWPOW)     return 32;
-   if (f == Algorithm::GHOSTRIDER) return 76;
-   return 39;
-}
-
 uint32_t xmrig::Job::getNumTransactions() const
 {
     if (!(m_algorithm.isCN() || m_algorithm.family() == Algorithm::RANDOM_X)) {

src/base/net/stratum/Job.h

@@ -7,8 +7,8 @@
  * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
  * Copyright 2018      Lee Clagett <https://github.com/vtnerd>
  * Copyright 2019      Howard Chu  <https://github.com/hyc>
- * Copyright 2018-2021 SChernykh   <https://github.com/SChernykh>
- * Copyright 2016-2021 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2024 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
  *
  *   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
@@ -27,11 +27,9 @@
 #ifndef XMRIG_JOB_H
 #define XMRIG_JOB_H
 
-
 #include <cstddef>
 #include <cstdint>
 
-
 #include "base/crypto/Algorithm.h"
 #include "base/tools/Buffer.h"
 #include "base/tools/String.h"
@@ -63,6 +61,7 @@ public:
     bool setBlob(const char *blob);
     bool setSeedHash(const char *hash);
     bool setTarget(const char *target);
+    size_t nonceOffset() const;
     void setDiff(uint64_t diff);
     void setSigKey(const char *sig_key);
 
@@ -77,7 +76,6 @@ public:
     inline const String &poolWallet() const             { return m_poolWallet; }
     inline const uint32_t *nonce() const                { return reinterpret_cast<const uint32_t*>(m_blob + nonceOffset()); }
     inline const uint8_t *blob() const                  { return m_blob; }
-    int32_t nonceOffset() const;
     inline size_t nonceSize() const                     { return (algorithm().family() == Algorithm::KAWPOW) ?  8 :  4; }
     inline size_t size() const                          { return m_size; }
     inline uint32_t *nonce()                            { return reinterpret_cast<uint32_t*>(m_blob + nonceOffset()); }
@@ -111,7 +109,7 @@ public:
 
     inline bool operator!=(const Job &other) const      { return !isEqual(other); }
     inline bool operator==(const Job &other) const      { return isEqual(other); }
-    inline Job &operator=(const Job &other)             { copy(other); return *this; }
+    inline Job &operator=(const Job &other)             { if (this != &other) { copy(other); } return *this; }
     inline Job &operator=(Job &&other) noexcept         { move(std::move(other)); return *this; }
 
 #   ifdef XMRIG_FEATURE_BENCHMARK

src/core/Miner.cpp

@@ -576,6 +576,11 @@ void xmrig::Miner::setJob(const Job &job, bool donate)
 
 #   ifdef XMRIG_ALGO_RANDOMX
     const bool ready = d_ptr->initRX();
+
+    // Always reset nonce on RandomX dataset change
+    if (!ready) {
+        d_ptr->reset = true;
+    }
 #   else
     constexpr const bool ready = true;
 #   endif

src/crypto/ghostrider/README.md

@@ -16,7 +16,7 @@ xmrig -a gr -o rtm.suprnova.cc:4273 --tls -u WALLET_ADDRESS -p x
 
 You can use **rtm_ghostrider_example.cmd** as a template and put pool URL and your wallet address there. The general XMRig documentation is available [here](https://xmrig.com/docs/miner).
 
-**Using `--threads` or `-t` option is NOT recommended because it turns off advanced built-in config.** If you want to tweak the nubmer of threads used for GhostRider, it's recommended to start using config.json instead of command line. The best suitable command line option for this is `--cpu-max-threads-hint=N` where N can be between 0 and 100.
+**Using `--threads` or `-t` option is NOT recommended because it turns off advanced built-in config.** If you want to tweak the number of threads used for GhostRider, it's recommended to start using config.json instead of command line. The best suitable command line option for this is `--cpu-max-threads-hint=N` where N can be between 0 and 100.
 
 ## Performance
 

src/crypto/randomx/jit_compiler_a64.cpp

@@ -94,7 +94,7 @@ static size_t CalcDatasetItemSize()
 		// Main loop prologue
 		((uint8_t*)randomx_calc_dataset_item_aarch64_mix - ((uint8_t*)randomx_calc_dataset_item_aarch64_prefetch)) + 4 +
 		// Inner main loop (instructions)
-		((RandomX_ConfigurationBase::SuperscalarLatency * 3) + 2) * 16 +
+		((RandomX_ConfigurationBase::SuperscalarMaxLatency * 3) + 2) * 16 +
 		// Main loop epilogue
 		((uint8_t*)randomx_calc_dataset_item_aarch64_store_result - (uint8_t*)randomx_calc_dataset_item_aarch64_mix) + 4
 	) +
@@ -1078,6 +1078,6 @@ void JitCompilerA64::h_NOP(Instruction& instr, uint32_t& codePos)
 {
 }
 
-InstructionGeneratorA64 JitCompilerA64::engine[257] = {};
+InstructionGeneratorA64 JitCompilerA64::engine[256] = {};
 
 }

src/crypto/randomx/jit_compiler_a64.hpp

@@ -74,7 +74,7 @@ namespace randomx {
 		void enableWriting() const;
 		void enableExecution() const;
 
-		static InstructionGeneratorA64 engine[257];
+		static InstructionGeneratorA64 engine[256];
 
 	private:
 		const bool hugePages;

src/crypto/randomx/jit_compiler_x86.cpp

@@ -266,6 +266,10 @@ namespace randomx {
 						// AVX2 init is slower on Zen4
 						initDatasetAVX2 = false;
 						break;
+					case xmrig::ICpuInfo::ARCH_ZEN5:
+						// AVX2 init is 49% faster on Zen5
+						initDatasetAVX2 = true;
+						break;
 					}
 				}
 			}
@@ -1443,6 +1447,6 @@ namespace randomx {
 		emitByte(0x90, code, codePos);
 	}
 
-	alignas(64) InstructionGeneratorX86 JitCompilerX86::engine[257] = {};
+	alignas(64) InstructionGeneratorX86 JitCompilerX86::engine[256] = {};
 
 }

src/crypto/randomx/jit_compiler_x86.hpp

@@ -81,7 +81,7 @@ namespace randomx {
 		void enableWriting() const;
 		void enableExecution() const;
 
-		alignas(64) static InstructionGeneratorX86 engine[257];
+		alignas(64) static InstructionGeneratorX86 engine[256];
 
 	private:
 		int registerUsage[RegistersCount] = {};

src/crypto/randomx/randomx.cpp

@@ -100,17 +100,18 @@ RandomX_ConfigurationSafex::RandomX_ConfigurationSafex()
 	ArgonSalt = "RandomSFX\x01";
 }
 
-RandomX_ConfigurationKeva::RandomX_ConfigurationKeva()
+RandomX_ConfigurationYada::RandomX_ConfigurationYada()
 {
-	ArgonSalt = "RandomKV\x01";
-	ScratchpadL2_Size = 131072;
-	ScratchpadL3_Size = 1048576;
+	ArgonSalt = "RandomXYadaCoin\x03";
+	SuperscalarLatency = 150;
+	ArgonIterations = 4;
 }
 
 RandomX_ConfigurationBase::RandomX_ConfigurationBase()
 	: ArgonIterations(3)
 	, ArgonLanes(1)
 	, ArgonSalt("RandomX\x03")
+	, SuperscalarLatency(170)
 	, ScratchpadL1_Size(16384)
 	, ScratchpadL2_Size(262144)
 	, ScratchpadL3_Size(2097152)
@@ -260,7 +261,7 @@ typedef void(randomx::JitCompilerX86::* InstructionGeneratorX86_2)(const randomx
 
 #define JIT_HANDLE(x, prev) do { \
 		const InstructionGeneratorX86_2 p = &randomx::JitCompilerX86::h_##x; \
-		memcpy(randomx::JitCompilerX86::engine + k, &p, sizeof(p)); \
+		memcpy(randomx::JitCompilerX86::engine + k, &p, sizeof(randomx::JitCompilerX86::engine[k])); \
 	} while (0)
 
 #elif (XMRIG_ARM == 8)
@@ -357,7 +358,7 @@ RandomX_ConfigurationWownero RandomX_WowneroConfig;
 RandomX_ConfigurationArqma RandomX_ArqmaConfig;
 RandomX_ConfigurationGraft RandomX_GraftConfig;
 RandomX_ConfigurationSafex RandomX_SafexConfig;
-RandomX_ConfigurationKeva RandomX_KevaConfig;
+RandomX_ConfigurationYada RandomX_YadaConfig;
 
 alignas(64) RandomX_ConfigurationBase RandomX_CurrentConfig;
 

src/crypto/randomx/randomx.h

@@ -69,7 +69,7 @@ struct RandomX_ConfigurationBase
 	{
 		ArgonMemory = 262144,
 		CacheAccesses = 8,
-		SuperscalarLatency = 170,
+		SuperscalarMaxLatency = 170,
 		DatasetBaseSize = 2147483648,
 		DatasetExtraSize = 33554368,
 		JumpBits = 8,
@@ -82,6 +82,7 @@ struct RandomX_ConfigurationBase
 	uint32_t ArgonIterations;
 	uint32_t ArgonLanes;
 	const char* ArgonSalt;
+	uint32_t SuperscalarLatency;
 
 	uint32_t ScratchpadL1_Size;
 	uint32_t ScratchpadL2_Size;
@@ -146,14 +147,14 @@ struct RandomX_ConfigurationWownero : public RandomX_ConfigurationBase { RandomX
 struct RandomX_ConfigurationArqma : public RandomX_ConfigurationBase { RandomX_ConfigurationArqma(); };
 struct RandomX_ConfigurationGraft : public RandomX_ConfigurationBase { RandomX_ConfigurationGraft(); };
 struct RandomX_ConfigurationSafex : public RandomX_ConfigurationBase { RandomX_ConfigurationSafex(); };
-struct RandomX_ConfigurationKeva : public RandomX_ConfigurationBase { RandomX_ConfigurationKeva(); };
+struct RandomX_ConfigurationYada : public RandomX_ConfigurationBase { RandomX_ConfigurationYada(); };
 
 extern RandomX_ConfigurationMonero RandomX_MoneroConfig;
 extern RandomX_ConfigurationWownero RandomX_WowneroConfig;
 extern RandomX_ConfigurationArqma RandomX_ArqmaConfig;
 extern RandomX_ConfigurationGraft RandomX_GraftConfig;
 extern RandomX_ConfigurationSafex RandomX_SafexConfig;
-extern RandomX_ConfigurationKeva RandomX_KevaConfig;
+extern RandomX_ConfigurationYada RandomX_YadaConfig;
 
 extern RandomX_ConfigurationBase RandomX_CurrentConfig;
 

src/crypto/rx/RxAlgo.cpp

@@ -44,8 +44,8 @@ const RandomX_ConfigurationBase *xmrig::RxAlgo::base(Algorithm::Id algorithm)
     case Algorithm::RX_SFX:
         return &RandomX_SafexConfig;
 
-    case Algorithm::RX_KEVA:
-        return &RandomX_KevaConfig;
+    case Algorithm::RX_YADA:
+        return &RandomX_YadaConfig;
 
     default:
         break;

src/crypto/rx/RxConfig.cpp

@@ -53,13 +53,15 @@ static const std::array<const char *, RxConfig::ModeMax> modeNames = { "auto", "
 
 
 #ifdef XMRIG_FEATURE_MSR
-constexpr size_t kMsrArraySize = 6;
+constexpr size_t kMsrArraySize = 7;
 
 static const std::array<MsrItems, kMsrArraySize> msrPresets = {
     MsrItems(),
     MsrItems{{ 0xC0011020, 0ULL }, { 0xC0011021, 0x40ULL, ~0x20ULL }, { 0xC0011022, 0x1510000ULL }, { 0xC001102b, 0x2000cc16ULL }},
     MsrItems{{ 0xC0011020, 0x0004480000000000ULL }, { 0xC0011021, 0x001c000200000040ULL, ~0x20ULL }, { 0xC0011022, 0xc000000401570000ULL }, { 0xC001102b, 0x2000cc10ULL }},
     MsrItems{{ 0xC0011020, 0x0004400000000000ULL }, { 0xC0011021, 0x0004000000000040ULL, ~0x20ULL }, { 0xC0011022, 0x8680000401570000ULL }, { 0xC001102b, 0x2040cc10ULL }},
+    MsrItems{{ 0xC0011020, 0x0004400000000000ULL }, { 0xC0011021, 0x0004000000000040ULL, ~0x20ULL }, { 0xC0011022, 0x8680000401570000ULL }, { 0xC001102b, 0x2040cc10ULL }},
+
     MsrItems{{ 0x1a4, 0xf }},
     MsrItems()
 };

src/version.h

@@ -22,15 +22,15 @@
 #define APP_ID        "xmrig"
 #define APP_NAME      "XMRig"
 #define APP_DESC      "XMRig miner"
-#define APP_VERSION   "6.21.2"
+#define APP_VERSION   "6.22.1-dev"
 #define APP_DOMAIN    "xmrig.com"
 #define APP_SITE      "www.xmrig.com"
 #define APP_COPYRIGHT "Copyright (C) 2016-2024 xmrig.com"
 #define APP_KIND      "miner"
 
 #define APP_VER_MAJOR  6
-#define APP_VER_MINOR  21
-#define APP_VER_PATCH  2
+#define APP_VER_MINOR  22
+#define APP_VER_PATCH  1
 
 #ifdef _MSC_VER
 #   if (_MSC_VER >= 1930)