Merge 753e63cd96 into a7be8cb80c

Update/Improve API docs
Cleanup API code
2025-12-25 21:52:54 -05:00 · 2024-06-20 08:25:13 +00:00 · 2023-07-12 02:06:53 -06:00 · 2023-07-12 02:06:53 -06:00
50 changed files with 4633 additions and 4472 deletions
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -1,11 +1,3 @@
-# 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.

--- a/doc/ALGORITHMS.md
+++ b/doc/ALGORITHMS.md
@@ -13,6 +13,7 @@ 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). |  |
--- a/doc/API.md
+++ b/doc/API.md
@@ -4,7 +4,7 @@ If you want use HTTP API you need enable it (`"enabled": true,`) then choice `po

 Offical HTTP client for API: http://workers.xmrig.info/

-Example configuration:
+Example configuration, used in Curl examples below:

 ```json
 "api": {
@@ -12,11 +12,11 @@ Example configuration:
 	"worker-id": null,
 },
 "http": {
-	"enabled": false,
+	"enabled": true,
 	"host": "127.0.0.1",
-	"port": 0,
-	"access-token": null,
-	"restricted": true
+	"port": 44444,
+	"access-token": "SECRET",
+	"restricted": false
 }
 ```

@@ -37,30 +37,78 @@ Versions before 2.15 was use another options for API https://github.com/xmrig/xm

 ## Endpoints

-### GET /1/summary
+### APIVersion 2

-Get miner summary information. [Example](api/1/summary.json).
+#### GET /2/summary

-### GET /1/threads
+Get miner summary information. [Example](api/2/summary.json).

-Get detailed information about miner threads. [Example](api/1/threads.json).
+#### GET /2/backends
+
+Get detailed information about miner backends. [Example](api/2/backends.json).
+
+### APIVersion 1 (deprecated)
+
+#### GET /1/summary
+
+Get miner summary information. Currently identical to `GET /2/summary`
+
+#### GET /1/threads
+
+**REMOVED** Get detailed information about miner threads. [Example](api/1/threads.json).
+
+Functionally replaced by `GET /2/backends` which contains a `threads` item per backend.
+
+### APIVersion 0 (deprecated)
+
+#### GET /api.json
+
+Get miner summary information. Currently identical to `GET /2/summary`


 ## Restricted endpoints

 All API endpoints below allow access to sensitive information and remote configure miner. You should set `access-token` and allow unrestricted access (`"restricted": false`).

-### GET /1/config
+### JSON-RPC Interface

-Get current miner configuration. [Example](api/1/config.json).
+#### POST /json_rpc

+Control miner with JSON-RPC. Methods: `pause`, `resume`, `stop`, `start`

-### PUT /1/config
+Curl example:
+
+```
+curl -v --data "{\"method\":\"pause\",\"id\":1}" -H "Content-Type: application/json" -H "Authorization: Bearer SECRET" http://127.0.0.1:44444/json_rpc
+```
+
+### APIVersion 2
+
+#### GET /2/config
+
+Get current miner configuration. [Example](api/2/config.json).
+
+#### PUT /2/config

 Update current miner configuration. Common use case, get current configuration, make changes, and upload it to miner.

 Curl example:

 ```
-curl -v --data-binary @config.json -X PUT -H "Content-Type: application/json" -H "Authorization: Bearer SECRET" http://127.0.0.1:44444/1/config
+...GET current config...
+curl -v -H "Content-Type: application/json" -H "Authorization: Bearer SECRET" http://127.0.0.1:44444/2/config > config.json
+...make changes...
+vim config.json
+...PUT changed config...
+curl -v --data-binary @config.json -X PUT -H "Content-Type: application/json" -H "Authorization: Bearer SECRET" http://127.0.0.1:44444/2/config
 ```
+
+### APIVersion 1 (deprecated)
+
+#### GET /1/config
+
+Get current miner configuration. Currently identical to `GET /2/config`
+
+#### PUT /1/config
+
+Update current miner configuration. Currently identical to `PUT /2/config`
--- a/doc/api/2/backends.json
+++ b/doc/api/2/backends.json
@@ -0,0 +1,78 @@
+[
+    {
+        "type": "cpu",
+        "enabled": true,
+        "algo": "rx/0",
+        "profile": "rx",
+        "hw-aes": true,
+        "priority": -1,
+        "msr": true,
+        "asm": "intel",
+        "argon2-impl": "AVX2",
+        "hugepages": [6, 6],
+        "memory": 6291456,
+        "hashrate": [1235.78, 1228.89, null],
+        "threads": [
+            {
+                "intensity": 1,
+                "affinity": 0,
+                "av": 1,
+                "hashrate": [409.75, 406.58, null]
+            },
+            {
+                "intensity": 1,
+                "affinity": 2,
+                "av": 1,
+                "hashrate": [412.9, 411.33, null]
+            },
+            {
+                "intensity": 1,
+                "affinity": 4,
+                "av": 1,
+                "hashrate": [413.11, 410.98, null]
+            }
+        ]
+    },
+    {
+        "type": "cuda",
+        "enabled": true,
+        "algo": "cn-heavy/xhv",
+        "profile": "cn-heavy/xhv",
+        "versions": {
+            "cuda-runtime": "11.6",
+            "cuda-driver": "11.6",
+            "plugin": "6.17.1-dev",
+            "nvml": "11.512.15",
+            "driver": "512.15"
+        },
+        "hashrate": [247.02, 247.34, null],
+        "threads": [
+            {
+                "index": 0,
+                "threads": 32,
+                "blocks": 38,
+                "bfactor": 6,
+                "bsleep": 25,
+                "affinity": -1,
+                "cclock": 0,
+                "mclock": 0,
+                "dataset_host": false,
+                "hashrate": [246.77, 247.26, null],
+                "name": "NVIDIA GeForce GTX 970",
+                "bus_id": "01:00.0",
+                "smx": 13,
+                "arch": 52,
+                "global_mem": 4294836224,
+                "clock": 1177,
+                "memory_clock": 3666,
+                "health": {
+                    "temperature": 69,
+                    "power": 161,
+                    "clock": 1328,
+                    "mem_clock": 3662,
+                    "fan_speed": [100]
+                }
+            }
+        ]
+    }
+]
--- a/doc/api/2/config.json
+++ b/doc/api/2/config.json
@@ -0,0 +1,136 @@
+{
+    "api": {
+        "id": null,
+        "worker-id": null
+    },
+    "http": {
+        "enabled": true,
+        "host": "127.0.0.1",
+        "port": 44444,
+        "access-token": "SECRET",
+        "restricted": false
+    },
+    "autosave": true,
+    "background": false,
+    "colors": true,
+    "title": true,
+    "randomx": {
+        "init": -1,
+        "init-avx2": -1,
+        "mode": "auto",
+        "1gb-pages": true,
+        "rdmsr": true,
+        "wrmsr": true,
+        "cache_qos": false,
+        "numa": true,
+        "scratchpad_prefetch_mode": 1
+    },
+    "cpu": {
+        "enabled": true,
+        "huge-pages": true,
+        "huge-pages-jit": true,
+        "hw-aes": null,
+        "priority": null,
+        "memory-pool": true,
+        "yield": true,
+        "asm": true,
+        "argon2-impl": null,
+        "argon2": [0, 2, 4, 6, 5, 7],
+        "astrobwt/v2": [1, 2, 3, 4, 5, 6, 7],
+        "cn": [
+            [1, 0],
+            [1, 2],
+            [1, 4]
+        ],
+        "cn-heavy": [
+            [1, 0],
+            [1, 2]
+        ],
+        "cn-lite": [
+            [1, 0],
+            [1, 2],
+            [1, 4],
+            [1, 6],
+            [1, 5],
+            [1, 7]
+        ],
+        "cn-pico": [
+            [2, 1],
+            [2, 2],
+            [2, 3],
+            [2, 4],
+            [2, 5],
+            [2, 6],
+            [2, 7]
+        ],
+        "cn/2": [
+            [1, 0],
+            [1, 2],
+            [1, 4]
+        ],
+        "cn/upx2": [
+            [2, 1],
+            [2, 2],
+            [2, 3],
+            [2, 4],
+            [2, 5],
+            [2, 6],
+            [2, 7]
+        ],
+        "ghostrider": [
+            [8, 0],
+            [8, 2],
+            [8, 4]
+        ],
+        "rx": [0, 2, 4],
+        "rx/arq": [1, 2, 3, 4, 5, 6, 7],
+        "rx/keva": [0, 2, 4, 6, 5, 7],
+        "rx/wow": [0, 2, 4, 6, 5, 7],
+        "cn-lite/0": false,
+        "cn/0": "cn"
+    },
+    "log-file": null,
+    "donate-level": 0,
+    "donate-over-proxy": 0,
+    "pools": [
+        {
+            "algo": null,
+            "coin": null,
+            "url": "some.pool:10064",
+            "user": "4blahblahblahblahblahblahblahblahblahblahblahblahblahblahblahblahblahblahblahblahblahblahblahbl",
+            "pass": "x",
+            "rig-id": null,
+            "nicehash": false,
+            "keepalive": true,
+            "enabled": true,
+            "tls": false,
+            "wss": false,
+            "daemon": false,
+            "socks5": null,
+            "self-select": null,
+            "submit-to-origin": false
+        }
+    ],
+    "retries": 5,
+    "retry-pause": 5,
+    "print-time": 64,
+    "syslog": false,
+    "tls": {
+        "enabled": false,
+        "protocols": null,
+        "cert": null,
+        "cert_key": null,
+        "ciphers": null,
+        "ciphersuites": null,
+        "dhparam": null
+    },
+    "dns": {
+        "ipv6": false,
+        "ttl": 30
+    },
+    "user-agent": null,
+    "verbose": 1,
+    "watch": true,
+    "pause-on-battery": false,
+    "pause-on-active": false
+}
--- a/doc/api/2/summary.json
+++ b/doc/api/2/summary.json
@@ -0,0 +1,76 @@
+{
+    "id": "51aca77da137cb62",
+    "worker_id": "tpad",
+    "uptime": 106,
+    "restricted": false,
+    "resources": {
+        "memory": {
+            "free": 4977348608,
+            "total": 16659546112,
+            "resident_set_memory": 16441344
+        },
+        "load_average": [3.4, 2.7, 2.44],
+        "hardware_concurrency": 8
+    },
+    "features": ["api", "asm", "http", "hwloc", "tls"],
+    "results": {
+        "diff_current": 37638,
+        "shares_good": 3,
+        "shares_total": 3,
+        "avg_time": 35,
+        "avg_time_ms": 35603,
+        "hashes_total": 165638,
+        "best": [358821, 344438, 73371, 0, 0, 0, 0, 0, 0, 0]
+    },
+    "algo": "rx/0",
+    "connection": {
+        "pool": "some.pool:20064",
+        "ip": "127.1.2.3",
+        "uptime": 106,
+        "uptime_ms": 106811,
+        "ping": 405,
+        "failures": 0,
+        "tls": "TLSv1.3",
+        "tls-fingerprint": null,
+        "algo": "rx/0",
+        "diff": 37638,
+        "accepted": 3,
+        "rejected": 0,
+        "avg_time": 35,
+        "avg_time_ms": 35603,
+        "hashes_total": 165638
+    },
+    "version": "6.17.1-dev",
+    "kind": "miner",
+    "ua": "XMRig/6.17.1-dev (Linux x86_64) libuv/1.43.0 gcc/9.4.0",
+    "cpu": {
+        "brand": "Intel(R) Core(TM) i7-4700MQ CPU @ 2.40GHz",
+        "family": 6,
+        "model": 60,
+        "stepping": 3,
+        "proc_info": 198339,
+        "aes": true,
+        "avx2": true,
+        "x64": true,
+        "64_bit": true,
+        "l2": 1048576,
+        "l3": 6291456,
+        "cores": 4,
+        "threads": 8,
+        "packages": 1,
+        "nodes": 1,
+        "backend": "hwloc/2.7.0",
+        "msr": "intel",
+        "assembly": "intel",
+        "arch": "x86_64",
+        "flags": ["aes", "avx", "avx2", "bmi2", "osxsave", "pdpe1gb", "sse2", "ssse3", "sse4.1", "popcnt"]
+    },
+    "donate_level": 0,
+    "paused": false,
+    "algorithms": ["cn/0", "cn/1", "cn/2", "cn/r", "cn/fast", "cn/half", "cn/xao", "cn/rto", "cn/rwz", "cn/zls", "cn/double", "cn/ccx", "cn-lite/1", "cn-heavy/0", "cn-heavy/tube", "cn-heavy/xhv", "cn-pico", "cn-pico/tlo", "cn/upx2", "rx/0", "rx/wow", "rx/arq", "rx/graft", "rx/sfx", "rx/keva", "argon2/chukwa", "argon2/chukwav2", "argon2/ninja", "astrobwt/v2", "ghostrider"],
+    "hashrate": {
+        "total": [1207.19, 1210.82, null],
+        "highest": 1316.85
+    },
+    "hugepages": [6, 6]
+}
--- a/scripts/build.openssl.sh
+++ b/scripts/build.openssl.sh
@@ -1,6 +1,6 @@
 #!/bin/sh -e

-OPENSSL_VERSION="1.1.1u"
+OPENSSL_VERSION="1.1.1s"

 mkdir -p deps
 mkdir -p deps/include
@@ -8,7 +8,7 @@ mkdir -p deps/lib

 mkdir -p build && cd build

-wget https://openssl.org/source/old/1.1.1/openssl-${OPENSSL_VERSION}.tar.gz -O openssl-${OPENSSL_VERSION}.tar.gz
+wget https://www.openssl.org/source/openssl-${OPENSSL_VERSION}.tar.gz -O openssl-${OPENSSL_VERSION}.tar.gz
 tar -xzf openssl-${OPENSSL_VERSION}.tar.gz

 cd openssl-${OPENSSL_VERSION}
--- a/scripts/build.openssl3.sh
+++ b/scripts/build.openssl3.sh
@@ -1,6 +1,6 @@
 #!/bin/sh -e

-OPENSSL_VERSION="3.0.14"
+OPENSSL_VERSION="3.0.13"

 mkdir -p deps
 mkdir -p deps/include
@@ -8,7 +8,7 @@ mkdir -p deps/lib

 mkdir -p build && cd build

-wget https://github.com/openssl/openssl/releases/download/openssl-${OPENSSL_VERSION}/openssl-${OPENSSL_VERSION}.tar.gz -O openssl-${OPENSSL_VERSION}.tar.gz
+wget https://www.openssl.org/source/openssl-${OPENSSL_VERSION}.tar.gz -O openssl-${OPENSSL_VERSION}.tar.gz
 tar -xzf openssl-${OPENSSL_VERSION}.tar.gz

 cd openssl-${OPENSSL_VERSION}
--- a/scripts/generate_cl.js
+++ b/scripts/generate_cl.js
@@ -50,6 +50,7 @@ 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',
--- a/src/backend/cpu/CpuConfig_gen.h
+++ b/src/backend/cpu/CpuConfig_gen.h
@@ -1,6 +1,6 @@
 /* XMRig
- * Copyright (c) 2018-2024 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2024 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright (c) 2018-2021 SChernykh   <https://github.com/SChernykh>
+ * Copyright (c) 2016-2021 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,6 +122,17 @@ 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));
    }
--- a/src/backend/cpu/interfaces/ICpuInfo.h
+++ b/src/backend/cpu/interfaces/ICpuInfo.h
@@ -52,8 +52,7 @@ public:
        ARCH_ZEN_PLUS,
        ARCH_ZEN2,
        ARCH_ZEN3,
-        ARCH_ZEN4,
-        ARCH_ZEN5
+        ARCH_ZEN4
    };

    enum MsrMod : uint32_t {
@@ -61,13 +60,12 @@ 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", "ryzen_1Ah_zen5", "intel", "custom"
+#   define MSR_NAMES_LIST "none", "ryzen_17h", "ryzen_19h", "ryzen_19h_zen4", "intel", "custom"

    enum Flag : uint32_t {
        FLAG_AES,
--- a/src/backend/cpu/platform/BasicCpuInfo.cpp
+++ b/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                                  = 7;
+constexpr size_t kMsrArraySize                                  = 6;
 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,11 +260,6 @@ 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;
--- a/src/backend/cuda/CudaConfig_gen.h
+++ b/src/backend/cuda/CudaConfig_gen.h
@@ -114,6 +114,7 @@ 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));
@@ -123,6 +124,10 @@ 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;
--- a/src/backend/opencl/OclWorker.cpp
+++ b/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()), m_job.nonceOffset(), results);
+                m_runner->run(readUnaligned(m_job.nonce()), results);
            }
            catch (std::exception &ex) {
                printError(id(), ex.what());
--- a/src/backend/opencl/cl/cn/algorithm.cl
+++ b/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
--- a/src/backend/opencl/cl/cn/cryptonight_cl.h
+++ b/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,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,
+    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,
    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,
--- a/src/backend/opencl/cl/rx/blake2b.cl
+++ b/src/backend/opencl/cl/rx/blake2b.cl
@@ -225,110 +225,6 @@ __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
--- a/src/backend/opencl/cl/rx/randomx.cl
+++ b/src/backend/opencl/cl/rx/randomx.cl
@@ -1,11 +1,13 @@
 #include "../cn/algorithm.cl"

-#if ((ALGO == ALGO_RX_0) || (ALGO == ALGO_RX_YADA))
+#if (ALGO == ALGO_RX_0)
 #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
--- a/src/backend/opencl/cl/rx/randomx_cl.h
+++ b/src/backend/opencl/cl/rx/randomx_cl.h
--- a/src/backend/opencl/cl/rx/randomx_constants_keva.h
+++ b/src/backend/opencl/cl/rx/randomx_constants_keva.h
@@ -0,0 +1,96 @@
+/*
+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)
--- a/src/backend/opencl/interfaces/IOclRunner.h
+++ b/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 nonce_offset, uint32_t *hashOutput)  = 0;
+    virtual void run(uint32_t nonce, uint32_t *hashOutput)  = 0;
    virtual void set(const Job &job, uint8_t *blob)         = 0;
    virtual void jobEarlyNotification(const Job&)           = 0;

--- a/src/backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.cpp
+++ b/src/backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.cpp
@@ -1,59 +0,0 @@
-/* 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);
-}
--- a/src/backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.h
+++ b/src/backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.h
@@ -1,50 +0,0 @@
-/* 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 */
--- a/src/backend/opencl/opencl.cmake
+++ b/src/backend/opencl/opencl.cmake
@@ -80,7 +80,6 @@ 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
@@ -98,7 +97,6 @@ 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
--- a/src/backend/opencl/runners/OclCnRunner.cpp
+++ b/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 /*nonce_offset*/, uint32_t *hashOutput)
+void xmrig::OclCnRunner::run(uint32_t nonce, uint32_t *hashOutput)
 {
    static const cl_uint zero = 0;

--- a/src/backend/opencl/runners/OclCnRunner.h
+++ b/src/backend/opencl/runners/OclCnRunner.h
@@ -42,7 +42,7 @@ public:

 protected:
    size_t bufferSize() const override;
-    void run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput) override;
+    void run(uint32_t nonce, uint32_t *hashOutput) override;
    void set(const Job &job, uint8_t *blob) override;
    void build() override;
    void init() override;
--- a/src/backend/opencl/runners/OclKawPowRunner.cpp
+++ b/src/backend/opencl/runners/OclKawPowRunner.cpp
@@ -75,7 +75,7 @@ OclKawPowRunner::~OclKawPowRunner()
 }


-void OclKawPowRunner::run(uint32_t nonce, uint32_t /*nonce_offset*/, uint32_t *hashOutput)
+void OclKawPowRunner::run(uint32_t nonce, uint32_t *hashOutput)
 {
    const size_t local_work_size = m_workGroupSize;
    const size_t global_work_offset = nonce;
--- a/src/backend/opencl/runners/OclKawPowRunner.h
+++ b/src/backend/opencl/runners/OclKawPowRunner.h
@@ -40,7 +40,7 @@ public:
    ~OclKawPowRunner() override;

 protected:
-    void run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput) override;
+    void run(uint32_t nonce, uint32_t *hashOutput) override;
    void set(const Job &job, uint8_t *blob) override;
    void build() override;
    void init() override;
--- a/src/backend/opencl/runners/OclRxBaseRunner.cpp
+++ b/src/backend/opencl/runners/OclRxBaseRunner.cpp
@@ -26,7 +26,6 @@
 #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"
@@ -74,7 +73,6 @@ 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;
@@ -87,7 +85,7 @@ xmrig::OclRxBaseRunner::~OclRxBaseRunner()
 }


-void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput)
+void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
 {
    static const uint32_t zero = 0;

@@ -98,7 +96,8 @@ void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t nonce_offset, uint32_t
        m_blake2b_initial_hash_double->setNonce(nonce);
    }
    else {
-        m_blake2b_initial_hash_big->setNonce(nonce, nonce_offset);
+        hashOutput[0xFF] = 0;
+        return;
    }

    m_find_shares->setNonce(nonce);
@@ -108,11 +107,8 @@ void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t nonce_offset, uint32_t
    if (m_jobSize <= 128) {
        m_blake2b_initial_hash->enqueue(m_queue, m_intensity);
    }
-    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_blake2b_initial_hash_double->enqueue(m_queue, m_intensity);
    }

    m_fillAes1Rx4_scratchpad->enqueue(m_queue, m_intensity);
@@ -154,15 +150,12 @@ 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());
 }
@@ -198,9 +191,6 @@ 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");

--- a/src/backend/opencl/runners/OclRxBaseRunner.h
+++ b/src/backend/opencl/runners/OclRxBaseRunner.h
@@ -36,7 +36,6 @@ namespace xmrig {
 class Blake2bHashRegistersKernel;
 class Blake2bInitialHashKernel;
 class Blake2bInitialHashDoubleKernel;
-class Blake2bInitialHashBigKernel;
 class FillAesKernel;
 class FindSharesKernel;
 class HashAesKernel;
@@ -54,7 +53,7 @@ protected:
    size_t bufferSize() const override;
    void build() override;
    void init() override;
-    void run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput) override;
+    void run(uint32_t nonce, uint32_t *hashOutput) override;
    void set(const Job &job, uint8_t *blob) override;

 protected:
@@ -64,7 +63,6 @@ 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;
--- a/src/base/api/interfaces/IApiRequest.h
+++ b/src/base/api/interfaces/IApiRequest.h
@@ -52,6 +52,8 @@ public:
    enum RequestType {
        REQ_UNKNOWN,
        REQ_SUMMARY,
+        REQ_BACKENDS,
+        REQ_CONFIG,
        REQ_JSON_RPC
    };

--- a/src/base/api/requests/ApiRequest.h
+++ b/src/base/api/requests/ApiRequest.h
@@ -60,7 +60,7 @@ protected:
    inline Source source() const override           { return m_source; }
    inline void done(int) override                  { m_state = STATE_DONE; }

-    int m_version       = 1;
+    int m_version       = 0;
    RequestType m_type  = REQ_UNKNOWN;
    State m_state       = STATE_NEW;
    String m_rpcMethod;
--- a/src/base/api/requests/HttpApiRequest.cpp
+++ b/src/base/api/requests/HttpApiRequest.cpp
@@ -67,10 +67,33 @@ xmrig::HttpApiRequest::HttpApiRequest(const HttpData &req, bool restricted) :
    m_res(req.id()),
    m_url(req.url.c_str())
 {
-    if (method() == METHOD_GET) {
-        if (url() == "/1/summary" || url() == "/2/summary" || url() == "/api.json") {
-            m_type = REQ_SUMMARY;
+    if (url().size() > 4 && memcmp(url().data(), "/", 1) == 0 && memcmp(url().data()+2, "/", 1) == 0) {
+        if (memcmp(url().data(), "/2/", 3) == 0) {
+            m_version = 2;
+        } else if (memcmp(url().data(), "/1/", 3) == 0) {
+            m_version = 1;
        }
+        switch (url().size()) {
+        case 9:
+            if (memcmp(url().data()+3, "config", 6) == 0) {
+                m_type = REQ_CONFIG;
+            }
+            break;
+        case 10:
+            if (memcmp(url().data()+3, "summary", 7) == 0) {
+                m_type = REQ_SUMMARY;
+            }
+            break;
+        case 11:
+            if (memcmp(url().data()+3, "backends", 8) == 0) {
+                m_type = REQ_BACKENDS;
+            }
+            break;
+        }
+    }
+
+    if (url() == "/api.json") {
+            m_type = REQ_SUMMARY;
    }

    if (method() == METHOD_POST && url() == "/json_rpc") {
@@ -94,12 +117,6 @@ xmrig::HttpApiRequest::HttpApiRequest(const HttpData &req, bool restricted) :

        return;
    }
-
-    if (url().size() > 4) {
-        if (memcmp(url().data(), "/2/", 3) == 0) {
-            m_version = 2;
-        }
-    }
 }


--- a/src/base/crypto/Algorithm.cpp
+++ b/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_YADA         = "rx/yada";
+const char *Algorithm::kRX_KEVA         = "rx/keva";
 #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_YADA),
+    ALGO_NAME(RX_KEVA),
 #   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_YADA),       ALGO_ALIAS(RX_YADA,         "randomx/yada"),
-                                    ALGO_ALIAS(RX_YADA,         "randomyada"),
+    ALGO_ALIAS_AUTO(RX_KEVA),       ALGO_ALIAS(RX_KEVA,         "randomx/keva"),
+                                    ALGO_ALIAS(RX_KEVA,         "randomkeva"),
 #   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_YADA,
+        RX_0, RX_WOW, RX_ARQ, RX_GRAFT, RX_SFX, RX_KEVA,
        AR2_CHUKWA, AR2_CHUKWA_V2, AR2_WRKZ,
        KAWPOW_RVN,
        GHOSTRIDER_RTM
--- a/src/base/crypto/Algorithm.h
+++ b/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_YADA         = 0x72151279,   // "rx/yada"          RandomYada (YadaCoin).
+        RX_KEVA         = 0x7214116b,   // "rx/keva"          RandomKEVA (Keva).
        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_YADA;
+    static const char *kRX_KEVA;
 #   endif

 #   ifdef XMRIG_ALGO_ARGON2
--- a/src/base/crypto/Coin.cpp
+++ b/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    ") },
 };


--- a/src/base/crypto/Coin.h
+++ b/src/base/crypto/Coin.h
@@ -36,11 +36,11 @@ public:
        SUMO,
        ARQMA,
        GRAFT,
+        KEVA,
        RAVEN,
        WOWNERO,
        ZEPHYR,
        TOWNFORGE,
-        YADA,
        MAX
    };

--- a/src/base/io/json/Json.cpp
+++ b/src/base/io/json/Json.cpp
@@ -211,13 +211,11 @@ rapidjson::Value xmrig::Json::normalize(double value, bool zero)
 {
    using namespace rapidjson;

-    const double value_rounded = floor(value * 100.0) / 100.0;
-
-    if (!std::isnormal(value) || !std::isnormal(value_rounded)) {
+    if (!std::isnormal(value)) {
        return zero ? Value(0.0) : Value(kNullType);
    }

-    return Value(value_rounded);
+    return Value(floor(value * 100.0) / 100.0);
 }


--- a/src/base/kernel/Base.cpp
+++ b/src/base/kernel/Base.cpp
@@ -45,13 +45,6 @@
 #ifdef XMRIG_FEATURE_API
 #   include "base/api/Api.h"
 #   include "base/api/interfaces/IApiRequest.h"
-
-namespace xmrig {
-
-static const char *kConfigPathV1 = "/1/config";
-static const char *kConfigPathV2 = "/2/config";
-
-} // namespace xmrig
 #endif


@@ -317,7 +310,7 @@ void xmrig::Base::onFileChanged(const String &fileName)
 void xmrig::Base::onRequest(IApiRequest &request)
 {
    if (request.method() == IApiRequest::METHOD_GET) {
-        if (request.url() == kConfigPathV1 || request.url() == kConfigPathV2) {
+        if (request.type() == IApiRequest::REQ_CONFIG) {
            if (request.isRestricted()) {
                return request.done(403);
            }
@@ -327,7 +320,7 @@ void xmrig::Base::onRequest(IApiRequest &request)
        }
    }
    else if (request.method() == IApiRequest::METHOD_PUT || request.method() == IApiRequest::METHOD_POST) {
-        if (request.url() == kConfigPathV1 || request.url() == kConfigPathV2) {
+        if (request.type() == IApiRequest::REQ_CONFIG) {
            request.accept();

            if (!reload(request.json())) {
--- a/src/base/net/stratum/Client.cpp
+++ b/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(Json::getString(params, "job_id"))) {
+    if (!job.setId(params["job_id"].GetString())) {
        *code = 3;
        return false;
    }
@@ -400,7 +400,7 @@ bool xmrig::Client::parseJob(const rapidjson::Value &params, int *code)
        }
    }

-    if (!job.setTarget(Json::getString(params, "target"))) {
+    if (!job.setTarget(params["target"].GetString())) {
        *code = 5;
        return false;
    }
--- a/src/base/net/stratum/Job.cpp
+++ b/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-2024 SChernykh   <https://github.com/SChernykh>
- * Copyright 2016-2024 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright 2018-2021 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2021 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,9 +24,11 @@
 *   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"
@@ -110,69 +112,35 @@ bool xmrig::Job::setSeedHash(const char *hash)

 bool xmrig::Job::setTarget(const char *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) {
+    if (!target) {
        return false;
    }

-    m_diff = toDiff(m_target);
+    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
-    if (size >= sizeof(m_rawTarget)) {
-        return false;
-    }
+    assert(sizeof(m_rawTarget) > (size * 2));

    memset(m_rawTarget, 0, sizeof(m_rawTarget));
-    memcpy(m_rawTarget, target, size);
+    memcpy(m_rawTarget, target, std::min(size * 2, sizeof(m_rawTarget)));
 #   endif

+    m_diff = toDiff(m_target);
    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;
@@ -203,6 +171,14 @@ 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)) {
--- a/src/base/net/stratum/Job.h
+++ b/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-2024 SChernykh   <https://github.com/SChernykh>
- * Copyright 2016-2024 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright 2018-2021 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2021 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,9 +27,11 @@
 #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"
@@ -61,7 +63,6 @@ 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);

@@ -76,6 +77,7 @@ 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()); }
@@ -109,7 +111,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)             { if (this != &other) { copy(other); } return *this; }
+    inline Job &operator=(const Job &other)             { copy(other); return *this; }
    inline Job &operator=(Job &&other) noexcept         { move(std::move(other)); return *this; }

 #   ifdef XMRIG_FEATURE_BENCHMARK
--- a/src/core/Miner.cpp
+++ b/src/core/Miner.cpp
@@ -689,7 +689,7 @@ void xmrig::Miner::onRequest(IApiRequest &request)
            d_ptr->getMiner(request.reply(), request.doc(), request.version());
            d_ptr->getHashrate(request.reply(), request.doc(), request.version());
        }
-        else if (request.url() == "/2/backends") {
+        else if (request.type() == IApiRequest::REQ_BACKENDS && request.version() == 2) {
            request.accept();

            d_ptr->getBackends(request.reply(), request.doc());
@@ -711,6 +711,12 @@ void xmrig::Miner::onRequest(IApiRequest &request)

            stop();
        }
+        else if (request.rpcMethod() == "start") {
+            request.accept();
+
+            const auto config = d_ptr->controller->config();
+            onConfigChanged(config, config);
+        }
    }

    for (IBackend *backend : d_ptr->backends) {
--- a/src/crypto/randomx/jit_compiler_a64.cpp
+++ b/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::SuperscalarMaxLatency * 3) + 2) * 16 +
+		((RandomX_ConfigurationBase::SuperscalarLatency * 3) + 2) * 16 +
 		// Main loop epilogue
 		((uint8_t*)randomx_calc_dataset_item_aarch64_store_result - (uint8_t*)randomx_calc_dataset_item_aarch64_mix) + 4
 	) +
--- a/src/crypto/randomx/jit_compiler_x86.cpp
+++ b/src/crypto/randomx/jit_compiler_x86.cpp
@@ -266,10 +266,6 @@ namespace randomx {
 						// AVX2 init is slower on Zen4
 						initDatasetAVX2 = false;
 						break;
-					case xmrig::ICpuInfo::ARCH_ZEN5:
-						// TODO: test it
-						initDatasetAVX2 = false;
-						break;
 					}
 				}
 			}
--- a/src/crypto/randomx/randomx.cpp
+++ b/src/crypto/randomx/randomx.cpp
@@ -100,18 +100,17 @@ RandomX_ConfigurationSafex::RandomX_ConfigurationSafex()
 	ArgonSalt = "RandomSFX\x01";
 }

-RandomX_ConfigurationYada::RandomX_ConfigurationYada()
+RandomX_ConfigurationKeva::RandomX_ConfigurationKeva()
 {
-	ArgonSalt = "RandomXYadaCoin\x03";
-	SuperscalarLatency = 150;
-	ArgonIterations = 4;
+	ArgonSalt = "RandomKV\x01";
+	ScratchpadL2_Size = 131072;
+	ScratchpadL3_Size = 1048576;
 }

 RandomX_ConfigurationBase::RandomX_ConfigurationBase()
 	: ArgonIterations(3)
 	, ArgonLanes(1)
 	, ArgonSalt("RandomX\x03")
-	, SuperscalarLatency(170)
 	, ScratchpadL1_Size(16384)
 	, ScratchpadL2_Size(262144)
 	, ScratchpadL3_Size(2097152)
@@ -358,7 +357,7 @@ RandomX_ConfigurationWownero RandomX_WowneroConfig;
 RandomX_ConfigurationArqma RandomX_ArqmaConfig;
 RandomX_ConfigurationGraft RandomX_GraftConfig;
 RandomX_ConfigurationSafex RandomX_SafexConfig;
-RandomX_ConfigurationYada RandomX_YadaConfig;
+RandomX_ConfigurationKeva RandomX_KevaConfig;

 alignas(64) RandomX_ConfigurationBase RandomX_CurrentConfig;

--- a/src/crypto/randomx/randomx.h
+++ b/src/crypto/randomx/randomx.h
@@ -69,7 +69,7 @@ struct RandomX_ConfigurationBase
 	{
 		ArgonMemory = 262144,
 		CacheAccesses = 8,
-		SuperscalarMaxLatency = 170,
+		SuperscalarLatency = 170,
 		DatasetBaseSize = 2147483648,
 		DatasetExtraSize = 33554368,
 		JumpBits = 8,
@@ -82,7 +82,6 @@ struct RandomX_ConfigurationBase
 	uint32_t ArgonIterations;
 	uint32_t ArgonLanes;
 	const char* ArgonSalt;
-	uint32_t SuperscalarLatency;

 	uint32_t ScratchpadL1_Size;
 	uint32_t ScratchpadL2_Size;
@@ -147,14 +146,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_ConfigurationYada : public RandomX_ConfigurationBase { RandomX_ConfigurationYada(); };
+struct RandomX_ConfigurationKeva : public RandomX_ConfigurationBase { RandomX_ConfigurationKeva(); };

 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_ConfigurationYada RandomX_YadaConfig;
+extern RandomX_ConfigurationKeva RandomX_KevaConfig;

 extern RandomX_ConfigurationBase RandomX_CurrentConfig;

--- a/src/crypto/rx/RxAlgo.cpp
+++ b/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_YADA:
-        return &RandomX_YadaConfig;
+    case Algorithm::RX_KEVA:
+        return &RandomX_KevaConfig;

    default:
        break;
--- a/src/crypto/rx/RxConfig.cpp
+++ b/src/crypto/rx/RxConfig.cpp
@@ -53,17 +53,13 @@ static const std::array<const char *, RxConfig::ModeMax> modeNames = { "auto", "


 #ifdef XMRIG_FEATURE_MSR
-constexpr size_t kMsrArraySize = 7;
+constexpr size_t kMsrArraySize = 6;

 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 }},
-
-    // TODO: Tune it for Zen5 when it's available
-    MsrItems{{ 0xC0011020, 0x0004400000000000ULL }, { 0xC0011021, 0x0004000000000040ULL, ~0x20ULL }, { 0xC0011022, 0x8680000401570000ULL }, { 0xC001102b, 0x2040cc10ULL }},
-
    MsrItems{{ 0x1a4, 0xf }},
    MsrItems()
 };
--- a/src/version.h
+++ b/src/version.h
@@ -22,15 +22,15 @@
 #define APP_ID        "xmrig"
 #define APP_NAME      "XMRig"
 #define APP_DESC      "XMRig miner"
-#define APP_VERSION   "6.22.0"
+#define APP_VERSION   "6.21.4-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  22
-#define APP_VER_PATCH  0
+#define APP_VER_MINOR  21
+#define APP_VER_PATCH  4

 #ifdef _MSC_VER
 #   if (_MSC_VER >= 1930)
Author	SHA1	Message	Date
Tony Butler	d6aa4a1591	Merge `753e63cd96` into `a7be8cb80c`	2024-06-20 08:25:13 +00:00
Tony Butler	753e63cd96	Update/Improve API docs	2023-07-12 02:06:53 -06:00
Tony Butler	f42b3e83a7	Cleanup API code	2023-07-12 02:06:53 -06:00