Update/Improve API docs

Cleanup API code
2025-12-08 00:15:04 -05:00 · 2023-07-12 02:06:53 -06:00 · 2023-07-12 02:06:53 -06:00
39 changed files with 3624 additions and 3863 deletions
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -1,18 +1,3 @@
-# v6.21.1
- [#3391](https://github.com/xmrig/xmrig/pull/3391) Added support for townforge (monero fork using randomx).
- [#3399](https://github.com/xmrig/xmrig/pull/3399) Fixed Zephyr mining (OpenCL).
- [#3420](https://github.com/xmrig/xmrig/pull/3420) Fixed segfault in HTTP API rebind.
-
-# v6.21.0
- [#3302](https://github.com/xmrig/xmrig/pull/3302) [#3312](https://github.com/xmrig/xmrig/pull/3312) Enabled keepalive for Windows (>= Vista).
- [#3320](https://github.com/xmrig/xmrig/pull/3320) Added "built for OS/architecture/bits" to "ABOUT".
- [#3339](https://github.com/xmrig/xmrig/pull/3339) Added SNI option for TLS connections.
- [#3342](https://github.com/xmrig/xmrig/pull/3342) Update `cn_main_loop.asm`.
- [#3346](https://github.com/xmrig/xmrig/pull/3346) ARM64 JIT: don't use `x18` register.
- [#3348](https://github.com/xmrig/xmrig/pull/3348) Update to latest `sse2neon.h`.
- [#3356](https://github.com/xmrig/xmrig/pull/3356) Updated pricing record size for **Zephyr** solo mining.
- [#3358](https://github.com/xmrig/xmrig/pull/3358) **Zephyr** solo mining: handle multiple outputs.
-
 # v6.20.0
 - Added new ARM CPU names.
 - [#2394](https://github.com/xmrig/xmrig/pull/2394) Added new CMake options `ARM_V8` and `ARM_V7`.
--- 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/src/backend/opencl/cl/rx/blake2b.cl
+++ b/src/backend/opencl/cl/rx/blake2b.cl
@@ -138,93 +138,6 @@ __kernel void blake2b_initial_hash(__global void *out, __global const void* bloc
 	t[7] = hash[7];
 }

-void blake2b_512_process_double_block_variable(ulong *out, ulong* m, __global const ulong* in, uint in_len, uint out_len)
-{
-	ulong v[16] =
-	{
-		iv0 ^ (0x01010000u | out_len), iv1, iv2, iv3, iv4      , iv5, iv6, iv7,
-		iv0               , iv1, iv2, iv3, iv4 ^ 128, iv5, iv6, iv7,
-	};
-
-	BLAKE2B_ROUNDS();
-
-	ulong h[8];
-	v[0] = h[0] = v[0] ^ v[8] ^ iv0 ^ (0x01010000u | out_len);
-	v[1] = h[1] = v[1] ^ v[9] ^ iv1;
-	v[2] = h[2] = v[2] ^ v[10] ^ iv2;
-	v[3] = h[3] = v[3] ^ v[11] ^ iv3;
-	v[4] = h[4] = v[4] ^ v[12] ^ iv4;
-	v[5] = h[5] = v[5] ^ v[13] ^ iv5;
-	v[6] = h[6] = v[6] ^ v[14] ^ iv6;
-	v[7] = h[7] = v[7] ^ v[15] ^ iv7;
-	v[8] = iv0;
-	v[9] = iv1;
-	v[10] = iv2;
-	v[11] = iv3;
-	v[12] = iv4 ^ in_len;
-	v[13] = iv5;
-	v[14] = ~iv6;
-	v[15] = iv7;
-
-	m[ 0] = (in_len > 128) ? in[16] : 0;
-	m[ 1] = (in_len > 136) ? in[17] : 0;
-	m[ 2] = (in_len > 144) ? in[18] : 0;
-	m[ 3] = (in_len > 152) ? in[19] : 0;
-	m[ 4] = (in_len > 160) ? in[20] : 0;
-	m[ 5] = (in_len > 168) ? in[21] : 0;
-	m[ 6] = (in_len > 176) ? in[22] : 0;
-	m[ 7] = (in_len > 184) ? in[23] : 0;
-	m[ 8] = (in_len > 192) ? in[24] : 0;
-	m[ 9] = (in_len > 200) ? in[25] : 0;
-	m[10] = (in_len > 208) ? in[26] : 0;
-	m[11] = (in_len > 216) ? in[27] : 0;
-	m[12] = (in_len > 224) ? in[28] : 0;
-	m[13] = (in_len > 232) ? in[29] : 0;
-	m[14] = (in_len > 240) ? in[30] : 0;
-	m[15] = (in_len > 248) ? in[31] : 0;
-
-	if (in_len % sizeof(ulong))
-		m[(in_len - 128) / sizeof(ulong)] &= (ulong)(-1) >> (64 - (in_len % sizeof(ulong)) * 8);
-
-	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_double(__global void *out, __global const void* blockTemplate, uint blockTemplateSize, uint start_nonce)
-{
-	const uint global_index = get_global_id(0);
-
-	__global const ulong* p = (__global const ulong*) blockTemplate;
-
-	ulong m[16] = { p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], p[8], p[9], p[10], p[11], p[12], p[13], p[14], p[15] };
-
-	const ulong nonce = start_nonce + global_index;
-	m[4] = (m[4] & ((ulong)(-1) >>  8)) | (nonce << 56);
-	m[5] = (m[5] & ((ulong)(-1) << 24)) | (nonce >>  8);
-
-	ulong hash[8];
-	blake2b_512_process_double_block_variable(hash, m, p, blockTemplateSize, 64);
-
-	__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.h
+++ b/src/backend/opencl/cl/rx/randomx_cl.h
--- a/src/backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.cpp
+++ b/src/backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.cpp
@@ -1,58 +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/Blake2bInitialHashDoubleKernel.h"
-#include "backend/opencl/wrappers/OclLib.h"
-
-
-void xmrig::Blake2bInitialHashDoubleKernel::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::Blake2bInitialHashDoubleKernel::setArgs(cl_mem out, cl_mem blockTemplate)
-{
-    setArg(0, sizeof(cl_mem), &out);
-    setArg(1, sizeof(cl_mem), &blockTemplate);
-}
-
-
-void xmrig::Blake2bInitialHashDoubleKernel::setBlobSize(size_t size)
-{
-    const uint32_t s = size;
-
-    setArg(2, sizeof(uint32_t), &s);
-}
-
-
-void xmrig::Blake2bInitialHashDoubleKernel::setNonce(uint32_t nonce)
-{
-    setArg(3, sizeof(uint32_t), &nonce);
-}
--- a/src/backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.h
+++ b/src/backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.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_BLAKE2BINITIALHASHDOUBLEKERNEL_H
-#define XMRIG_BLAKE2BINITIALHASHDOUBLEKERNEL_H
-
-
-#include "backend/opencl/wrappers/OclKernel.h"
-
-
-namespace xmrig {
-
-
-class Blake2bInitialHashDoubleKernel : public OclKernel
-{
-public:
-    inline Blake2bInitialHashDoubleKernel(cl_program program) : OclKernel(program, "blake2b_initial_hash_double") {}
-
-    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);
-};
-
-
-} // namespace xmrig
-
-
-#endif /* XMRIG_BLAKE2BINITIALHASHDOUBLEKERNEL_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/Blake2bInitialHashDoubleKernel.h
             src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.h
             src/backend/opencl/kernels/rx/ExecuteVmKernel.h
             src/backend/opencl/kernels/rx/FillAesKernel.h
@@ -97,7 +96,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/Blake2bInitialHashDoubleKernel.cpp
             src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.cpp
             src/backend/opencl/kernels/rx/ExecuteVmKernel.cpp
             src/backend/opencl/kernels/rx/FillAesKernel.cpp
--- a/src/backend/opencl/runners/OclRxBaseRunner.cpp
+++ b/src/backend/opencl/runners/OclRxBaseRunner.cpp
@@ -25,7 +25,6 @@
 #include "backend/opencl/runners/OclRxBaseRunner.h"
 #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/FillAesKernel.h"
 #include "backend/opencl/kernels/rx/FindSharesKernel.h"
 #include "backend/opencl/kernels/rx/HashAesKernel.h"
@@ -72,7 +71,6 @@ xmrig::OclRxBaseRunner::~OclRxBaseRunner()
    delete m_fillAes4Rx4_entropy;
    delete m_hashAes1Rx4;
    delete m_blake2b_initial_hash;
-    delete m_blake2b_initial_hash_double;
    delete m_blake2b_hash_registers_32;
    delete m_blake2b_hash_registers_64;
    delete m_find_shares;
@@ -89,28 +87,12 @@ void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
 {
    static const uint32_t zero = 0;

-    if (m_jobSize <= 128) {
-        m_blake2b_initial_hash->setNonce(nonce);
-    }
-    else if (m_jobSize <= 256) {
-        m_blake2b_initial_hash_double->setNonce(nonce);
-    }
-    else {
-        hashOutput[0xFF] = 0;
-        return;
-    }
-
+    m_blake2b_initial_hash->setNonce(nonce);
    m_find_shares->setNonce(nonce);

    enqueueWriteBuffer(m_output, CL_FALSE, sizeof(cl_uint) * 0xFF, sizeof(uint32_t), &zero);

-    if (m_jobSize <= 128) {
-        m_blake2b_initial_hash->enqueue(m_queue, m_intensity);
-    }
-    else {
-        m_blake2b_initial_hash_double->enqueue(m_queue, m_intensity);
-    }
-
+    m_blake2b_initial_hash->enqueue(m_queue, m_intensity);
    m_fillAes1Rx4_scratchpad->enqueue(m_queue, m_intensity);

    const uint32_t programCount = RxAlgo::programCount(m_algorithm);
@@ -152,11 +134,7 @@ void xmrig::OclRxBaseRunner::set(const Job &job, uint8_t *blob)

    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_find_shares->setTarget(job.target());
 }

@@ -188,9 +166,6 @@ void xmrig::OclRxBaseRunner::build()
    m_blake2b_initial_hash = new Blake2bInitialHashKernel(m_program);
    m_blake2b_initial_hash->setArgs(m_hashes, m_input);

-    m_blake2b_initial_hash_double = new Blake2bInitialHashDoubleKernel(m_program);
-    m_blake2b_initial_hash_double->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
@@ -35,7 +35,6 @@ namespace xmrig {

 class Blake2bHashRegistersKernel;
 class Blake2bInitialHashKernel;
-class Blake2bInitialHashDoubleKernel;
 class FillAesKernel;
 class FindSharesKernel;
 class HashAesKernel;
@@ -59,24 +58,21 @@ protected:
 protected:
    virtual void execute(uint32_t iteration) = 0;

-    Blake2bHashRegistersKernel *m_blake2b_hash_registers_32       = nullptr;
-    Blake2bHashRegistersKernel *m_blake2b_hash_registers_64       = nullptr;
-    Blake2bInitialHashKernel *m_blake2b_initial_hash              = nullptr;
-    Blake2bInitialHashDoubleKernel *m_blake2b_initial_hash_double = nullptr;
+    Blake2bHashRegistersKernel *m_blake2b_hash_registers_32 = nullptr;
+    Blake2bHashRegistersKernel *m_blake2b_hash_registers_64 = nullptr;
+    Blake2bInitialHashKernel *m_blake2b_initial_hash        = nullptr;
    Buffer m_seed;
-    cl_mem m_dataset                                              = nullptr;
-    cl_mem m_entropy                                              = nullptr;
-    cl_mem m_hashes                                               = nullptr;
-    cl_mem m_rounding                                             = nullptr;
-    cl_mem m_scratchpads                                          = nullptr;
-    FillAesKernel *m_fillAes1Rx4_scratchpad                       = nullptr;
-    FillAesKernel *m_fillAes4Rx4_entropy                          = nullptr;
-    FindSharesKernel *m_find_shares                               = nullptr;
-    HashAesKernel *m_hashAes1Rx4                                  = nullptr;
-    uint32_t m_gcn_version                                        = 12;
-    uint32_t m_worksize                                           = 8;
-
-    size_t m_jobSize                                              = 0;
+    cl_mem m_dataset                                        = nullptr;
+    cl_mem m_entropy                                        = nullptr;
+    cl_mem m_hashes                                         = nullptr;
+    cl_mem m_rounding                                       = nullptr;
+    cl_mem m_scratchpads                                    = nullptr;
+    FillAesKernel *m_fillAes1Rx4_scratchpad                 = nullptr;
+    FillAesKernel *m_fillAes4Rx4_entropy                    = nullptr;
+    FindSharesKernel *m_find_shares                         = nullptr;
+    HashAesKernel *m_hashAes1Rx4                            = nullptr;
+    uint32_t m_gcn_version                                  = 12;
+    uint32_t m_worksize                                     = 8;
 };


--- a/src/base/api/Api.cpp
+++ b/src/base/api/Api.cpp
@@ -39,7 +39,6 @@


 #include <thread>
-#include <iostream>


 namespace xmrig {
@@ -81,8 +80,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);

@@ -93,11 +91,7 @@ xmrig::Api::Api(Base *base) :
 xmrig::Api::~Api()
 {
 #   ifdef XMRIG_FEATURE_HTTP
-    if (m_httpd) {
-        m_httpd->stop();
-        delete m_httpd;
-        m_httpd = nullptr; // Ensure the pointer is set to nullptr after deletion
-    }
+    delete m_httpd;
 #   endif
 }

@@ -115,14 +109,8 @@ void xmrig::Api::start()
    genWorkerId(m_base->config()->apiWorkerId());

 #   ifdef XMRIG_FEATURE_HTTP
-    if (!m_httpd) {
-        m_httpd = new Httpd(m_base);
-        if (!m_httpd->start()) {
-            std::cerr << "HTTP server failed to start." << std::endl;
-            delete m_httpd; // Properly handle failure to start
-            m_httpd = nullptr;
-        }
-    }
+    m_httpd = new Httpd(m_base);
+    m_httpd->start();
 #   endif
 }

@@ -130,9 +118,7 @@ void xmrig::Api::start()
 void xmrig::Api::stop()
 {
 #   ifdef XMRIG_FEATURE_HTTP
-    if (m_httpd) {
-        m_httpd->stop();
-    }
+    m_httpd->stop();
 #   endif
 }

@@ -140,15 +126,13 @@ void xmrig::Api::stop()
 void xmrig::Api::tick()
 {
 #   ifdef XMRIG_FEATURE_HTTP
-    if (!m_httpd || !m_base->config()->http().isEnabled() || m_httpd->isBound()) {
+    if (m_httpd->isBound() || !m_base->config()->http().isEnabled()) {
        return;
    }

    if (++m_ticks % 10 == 0) {
        m_ticks = 0;
-        if (m_httpd) {
-            m_httpd->start();
-        }
+        m_httpd->start();
    }
 #   endif
 }
--- 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/Coin.cpp
+++ b/src/base/crypto/Coin.cpp
@@ -54,7 +54,6 @@ static const CoinInfo coinInfo[] = {
    { 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 ") },
 };


--- a/src/base/crypto/Coin.h
+++ b/src/base/crypto/Coin.h
@@ -40,7 +40,6 @@ public:
        RAVEN,
        WOWNERO,
        ZEPHYR,
-        TOWNFORGE,
        MAX
    };

--- 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/kernel/Platform.h
+++ b/src/base/kernel/Platform.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-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
@@ -49,12 +49,6 @@ public:
    static inline bool isUserActive(uint64_t ms)    { return idleTime() < ms; }
    static inline const String &userAgent()         { return m_userAgent; }

-#   ifdef XMRIG_OS_WIN
-    static bool hasKeepalive();
-#   else
-    static constexpr bool hasKeepalive()            { return true; }
-#   endif
-
    static bool isOnBatteryPower();
    static uint64_t idleTime();

--- a/src/base/kernel/Platform_win.cpp
+++ b/src/base/kernel/Platform_win.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-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
@@ -16,6 +16,7 @@
 *   along with this program. If not, see <http://www.gnu.org/licenses/>.
 */

+
 #include <algorithm>
 #include <winsock2.h>
 #include <windows.h>
@@ -69,12 +70,6 @@ char *xmrig::Platform::createUserAgent()
 }


-bool xmrig::Platform::hasKeepalive()
-{
-    return winOsVersion().dwMajorVersion >= 6;
-}
-
-
 #ifndef XMRIG_FEATURE_HWLOC
 bool xmrig::Platform::setThreadAffinity(uint64_t cpu_id)
 {
--- a/src/base/kernel/config/BaseConfig.cpp
+++ b/src/base/kernel/config/BaseConfig.cpp
@@ -142,7 +142,7 @@ void xmrig::BaseConfig::printVersions()
    snprintf(buf, sizeof buf, "MSVC/%d", MSVC_VERSION);
 #   endif

-    Log::print(GREEN_BOLD(" * ") WHITE_BOLD("%-13s") CYAN_BOLD("%s/%s") WHITE_BOLD(" %s") WHITE_BOLD(" (built for %s") WHITE_BOLD(" %s,") WHITE_BOLD(" %s)"), "ABOUT", APP_NAME, APP_VERSION, buf, APP_OS, APP_ARCH, APP_BITS);
+    Log::print(GREEN_BOLD(" * ") WHITE_BOLD("%-13s") CYAN_BOLD("%s/%s") WHITE_BOLD(" %s"), "ABOUT", APP_NAME, APP_VERSION, buf);

    std::string libs;

--- a/src/base/net/stratum/Client.cpp
+++ b/src/base/net/stratum/Client.cpp
@@ -1,7 +1,7 @@
 /* XMRig
 * Copyright (c) 2019      jtgrassie   <https://github.com/jtgrassie>
- * Copyright (c) 2018-2023 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2023 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
@@ -42,14 +42,13 @@
 #include "base/io/json/JsonRequest.h"
 #include "base/io/log/Log.h"
 #include "base/kernel/interfaces/IClientListener.h"
-#include "base/kernel/Platform.h"
 #include "base/net/dns/Dns.h"
 #include "base/net/dns/DnsRecords.h"
 #include "base/net/stratum/Socks5.h"
 #include "base/net/tools/NetBuffer.h"
 #include "base/tools/Chrono.h"
-#include "base/tools/cryptonote/BlobReader.h"
 #include "base/tools/Cvt.h"
+#include "base/tools/cryptonote/BlobReader.h"
 #include "net/JobResult.h"


@@ -344,9 +343,6 @@ bool xmrig::Client::close()
    setState(ClosingState);

    if (uv_is_closing(reinterpret_cast<uv_handle_t*>(m_socket)) == 0) {
-        if (Platform::hasKeepalive()) {
-            uv_tcp_keepalive(m_socket, 0, 60);
-        }
        uv_close(reinterpret_cast<uv_handle_t*>(m_socket), Client::onClose);
    }

@@ -571,9 +567,9 @@ void xmrig::Client::connect(const sockaddr *addr)
    uv_tcp_init(uv_default_loop(), m_socket);
    uv_tcp_nodelay(m_socket, 1);

-    if (Platform::hasKeepalive()) {
-        uv_tcp_keepalive(m_socket, 1, 60);
-    }
+#   ifndef WIN32
+    uv_tcp_keepalive(m_socket, 1, 60);
+#   endif

    uv_tcp_connect(req, m_socket, addr, onConnect);
 }
@@ -589,7 +585,7 @@ void xmrig::Client::handshake()
    if (isTLS()) {
        m_expire = Chrono::steadyMSecs() + kResponseTimeout;

-        m_tls->handshake(m_pool.isSNI() ? m_pool.host().data() : nullptr);
+        m_tls->handshake();
    }
    else
 #   endif
--- a/src/base/net/stratum/DaemonClient.cpp
+++ b/src/base/net/stratum/DaemonClient.cpp
@@ -1,13 +1,13 @@
 /* XMRig
- * Copyright (c) 2010      Jeff Garzik <jgarzik@pobox.com>
- * Copyright (c) 2012-2014 pooler      <pooler@litecoinpool.org>
- * Copyright (c) 2014      Lucas Jones <https://github.com/lucasjones>
- * Copyright (c) 2014-2016 Wolf9466    <https://github.com/OhGodAPet>
- * Copyright (c) 2016      Jay D Dee   <jayddee246@gmail.com>
- * Copyright (c) 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
- * Copyright (c) 2019      Howard Chu  <https://github.com/hyc>
- * Copyright (c) 2018-2023 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2023 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * 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 2019      Howard Chu  <https://github.com/hyc>
+ * Copyright 2018-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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,6 +23,7 @@
 *   along with this program. If not, see <http://www.gnu.org/licenses/>.
 */

+
 #include <uv.h>


@@ -33,7 +34,6 @@
 #include "base/io/json/JsonRequest.h"
 #include "base/io/log/Log.h"
 #include "base/kernel/interfaces/IClientListener.h"
-#include "base/kernel/Platform.h"
 #include "base/net/dns/Dns.h"
 #include "base/net/dns/DnsRecords.h"
 #include "base/net/http/Fetch.h"
@@ -42,9 +42,9 @@
 #include "base/net/stratum/SubmitResult.h"
 #include "base/net/tools/NetBuffer.h"
 #include "base/tools/bswap_64.h"
-#include "base/tools/cryptonote/Signatures.h"
 #include "base/tools/Cvt.h"
 #include "base/tools/Timer.h"
+#include "base/tools/cryptonote/Signatures.h"
 #include "net/JobResult.h"


@@ -358,9 +358,9 @@ void xmrig::DaemonClient::onResolved(const DnsRecords &records, int status, cons
    uv_tcp_init(uv_default_loop(), s);
    uv_tcp_nodelay(s, 1);

-    if (Platform::hasKeepalive()) {
-        uv_tcp_keepalive(s, 1, 60);
-    }
+#   ifndef WIN32
+    uv_tcp_keepalive(s, 1, 60);
+#   endif

    if (m_pool.zmq_port() > 0) {
        delete m_ZMQSocket;
@@ -589,9 +589,6 @@ void xmrig::DaemonClient::retry()
    }

    if ((m_ZMQConnectionState != ZMQ_NOT_CONNECTED) && (m_ZMQConnectionState != ZMQ_DISCONNECTING)) {
-        if (Platform::hasKeepalive()) {
-            uv_tcp_keepalive(m_ZMQSocket, 0, 60);
-        }
        uv_close(reinterpret_cast<uv_handle_t*>(m_ZMQSocket), onZMQClose);
    }

@@ -919,9 +916,6 @@ bool xmrig::DaemonClient::ZMQClose(bool shutdown)
    m_ZMQConnectionState = ZMQ_DISCONNECTING;

    if (uv_is_closing(reinterpret_cast<uv_handle_t*>(m_ZMQSocket)) == 0) {
-        if (Platform::hasKeepalive()) {
-            uv_tcp_keepalive(m_ZMQSocket, 0, 60);
-        }
        uv_close(reinterpret_cast<uv_handle_t*>(m_ZMQSocket), shutdown ? onZMQShutdown : onZMQClose);
        if (!shutdown) {
            retry();
--- a/src/base/net/stratum/Pool.cpp
+++ b/src/base/net/stratum/Pool.cpp
@@ -77,7 +77,6 @@ const char *Pool::kSelfSelect             = "self-select";
 const char *Pool::kSOCKS5                 = "socks5";
 const char *Pool::kSubmitToOrigin         = "submit-to-origin";
 const char *Pool::kTls                    = "tls";
-const char *Pool::kSni                    = "sni";
 const char *Pool::kUrl                    = "url";
 const char *Pool::kUser                   = "user";
 const char *Pool::kSpendSecretKey         = "spend-secret-key";
@@ -138,7 +137,6 @@ xmrig::Pool::Pool(const rapidjson::Value &object) :
    m_flags.set(FLAG_ENABLED,  Json::getBool(object, kEnabled, true));
    m_flags.set(FLAG_NICEHASH, Json::getBool(object, kNicehash) || m_url.host().contains(kNicehashHost));
    m_flags.set(FLAG_TLS,      Json::getBool(object, kTls) || m_url.isTLS());
-    m_flags.set(FLAG_SNI,      Json::getBool(object, kSni));

    setKeepAlive(Json::getValue(object, kKeepalive));

@@ -301,7 +299,6 @@ rapidjson::Value xmrig::Pool::toJSON(rapidjson::Document &doc) const

    obj.AddMember(StringRef(kEnabled),      m_flags.test(FLAG_ENABLED), allocator);
    obj.AddMember(StringRef(kTls),          isTLS(), allocator);
-    obj.AddMember(StringRef(kSni),          isSNI(), allocator);
    obj.AddMember(StringRef(kFingerprint),  m_fingerprint.toJSON(), allocator);
    obj.AddMember(StringRef(kDaemon),       m_mode == MODE_DAEMON, allocator);
    obj.AddMember(StringRef(kSOCKS5),       m_proxy.toJSON(doc), allocator);
--- a/src/base/net/stratum/Pool.h
+++ b/src/base/net/stratum/Pool.h
@@ -1,7 +1,7 @@
 /* XMRig
 * Copyright (c) 2019      Howard Chu  <https://github.com/hyc>
- * 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
@@ -59,7 +59,7 @@ public:
    static const char *kCoin;
    static const char *kDaemon;
    static const char *kDaemonPollInterval;
-    static const char *kDaemonJobTimeout;
+    static const char* kDaemonJobTimeout;
    static const char *kEnabled;
    static const char *kFingerprint;
    static const char *kKeepalive;
@@ -70,11 +70,10 @@ public:
    static const char *kSOCKS5;
    static const char *kSubmitToOrigin;
    static const char *kTls;
-    static const char *kSni;
    static const char *kUrl;
    static const char *kUser;
-    static const char *kSpendSecretKey;
-    static const char *kDaemonZMQPort;
+    static const char* kSpendSecretKey;
+    static const char* kDaemonZMQPort;
    static const char *kNicehashHost;

    constexpr static int kKeepAliveTimeout         = 60;
@@ -96,7 +95,6 @@ public:

    inline bool isNicehash() const                      { return m_flags.test(FLAG_NICEHASH); }
    inline bool isTLS() const                           { return m_flags.test(FLAG_TLS) || m_url.isTLS(); }
-    inline bool isSNI() const                           { return m_flags.test(FLAG_SNI); }
    inline bool isValid() const                         { return m_url.isValid(); }
    inline const Algorithm &algorithm() const           { return m_algorithm; }
    inline const Coin &coin() const                     { return m_coin; }
@@ -140,7 +138,6 @@ private:
        FLAG_ENABLED,
        FLAG_NICEHASH,
        FLAG_TLS,
-        FLAG_SNI,
        FLAG_MAX
    };

--- a/src/base/net/stratum/Tls.cpp
+++ b/src/base/net/stratum/Tls.cpp
@@ -60,7 +60,7 @@ xmrig::Client::Tls::~Tls()
 }


-bool xmrig::Client::Tls::handshake(const char* servername)
+bool xmrig::Client::Tls::handshake()
 {
    m_ssl = SSL_new(m_ctx);
    assert(m_ssl != nullptr);
@@ -69,10 +69,6 @@ bool xmrig::Client::Tls::handshake(const char* servername)
        return false;
    }

-    if (servername) {
-        SSL_set_tlsext_host_name(m_ssl, servername);
-    }
-
    SSL_set_connect_state(m_ssl);
    SSL_set_bio(m_ssl, m_read, m_write);
    SSL_do_handshake(m_ssl);
--- a/src/base/net/stratum/Tls.h
+++ b/src/base/net/stratum/Tls.h
@@ -42,7 +42,7 @@ public:
    Tls(Client *client);
    ~Tls();

-    bool handshake(const char* servername);
+    bool handshake();
    bool send(const char *data, size_t size);
    const char *fingerprint() const;
    const char *version() const;
--- a/src/base/tools/cryptonote/BlockTemplate.cpp
+++ b/src/base/tools/cryptonote/BlockTemplate.cpp
@@ -1,8 +1,8 @@
 /* XMRig
 * Copyright (c) 2012-2013 The Cryptonote developers
 * Copyright (c) 2014-2021 The Monero Project
- * Copyright (c) 2018-2023 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2023 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
@@ -198,7 +198,7 @@ bool xmrig::BlockTemplate::parse(bool hashes)
    }

    if (m_coin == Coin::ZEPHYR) {
-        uint8_t pricing_record[120];
+        uint8_t pricing_record[24];
        ar(pricing_record);
    }

@@ -207,11 +207,7 @@ bool xmrig::BlockTemplate::parse(bool hashes)
    setOffset(MINER_TX_PREFIX_OFFSET, ar.index());

    ar(m_txVersion);
-
-    if (m_coin != Coin::TOWNFORGE) {
-      ar(m_unlockTime);
-    }
-
+    ar(m_unlockTime);
    ar(m_numInputs);

    // must be 1 input
@@ -229,12 +225,8 @@ bool xmrig::BlockTemplate::parse(bool hashes)
    ar(m_height);
    ar(m_numOutputs);

-    if (m_coin == Coin::ZEPHYR) {
-        if (m_numOutputs < 2) {
-            return false;
-        }
-    }
-    else if (m_numOutputs != 1) {
+    const uint64_t expected_outputs = (m_coin == Coin::ZEPHYR) ? 2 : 1;
+    if (m_numOutputs != expected_outputs) {
        return false;
    }

@@ -260,34 +252,28 @@ bool xmrig::BlockTemplate::parse(bool hashes)
        ar.skip(asset_type_len);
        ar(m_viewTag);

-        for (uint64_t k = 1; k < m_numOutputs; ++k) {
-            uint64_t amount2;
-            ar(amount2);
+        uint64_t amount2;
+        ar(amount2);

-            uint8_t output_type2;
-            ar(output_type2);
-            if (output_type2 != 2) {
-                return false;
-            }
-
-            Span key2;
-            ar(key2, kKeySize);
-
-            ar(asset_type_len);
-            ar.skip(asset_type_len);
-
-            uint8_t view_tag2;
-            ar(view_tag2);
+        uint8_t output_type2;
+        ar(output_type2);
+        if (output_type2 != 2) {
+            return false;
        }
+
+        Span key2;
+        ar(key2, kKeySize);
+
+        ar(asset_type_len);
+        ar.skip(asset_type_len);
+
+        uint8_t view_tag2;
+        ar(view_tag2);
    }
    else if (m_outputType == 3) {
        ar(m_viewTag);
    }

-    if (m_coin == Coin::TOWNFORGE) {
-      ar(m_unlockTime);
-    }
-
    ar(m_extraSize);

    setOffset(TX_EXTRA_OFFSET, ar.index());
@@ -343,11 +329,6 @@ bool xmrig::BlockTemplate::parse(bool hashes)
    uint8_t vin_rct_type = 0;
    ar(vin_rct_type);

-    // no way I'm parsing a full game update here
-    if (m_coin == Coin::TOWNFORGE && m_height % 720 == 0) {
-      return true;
-    }
-
    // must be RCTTypeNull (0)
    if (vin_rct_type != 0) {
        return false;
--- a/src/base/tools/cryptonote/WalletAddress.cpp
+++ b/src/base/tools/cryptonote/WalletAddress.cpp
@@ -1,8 +1,8 @@
 /* XMRig
 * Copyright (c) 2012-2013 The Cryptonote developers
 * Copyright (c) 2014-2021 The Monero Project
- * Copyright (c) 2018-2023 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2023 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
@@ -33,25 +33,6 @@

 bool xmrig::WalletAddress::decode(const char *address, size_t size)
 {
-    uint64_t tf_tag = 0;
-    if (size >= 4 && !strncmp(address, "TF", 2)) {
-      tf_tag = 0x424200;
-      switch (address[2])
-      {
-        case '1': tf_tag |= 0; break;
-        case '2': tf_tag |= 1; break;
-        default: tf_tag = 0; return false;
-      }
-      switch (address[3]) {
-        case 'M': tf_tag |= 0; break;
-        case 'T': tf_tag |= 0x10; break;
-        case 'S': tf_tag |= 0x20; break;
-        default: tf_tag = 0; return false;
-      }
-      address += 4;
-      size -= 4;
-    }
-
    static constexpr std::array<int, 9> block_sizes{ 0, 2, 3, 5, 6, 7, 9, 10, 11 };
    static constexpr char alphabet[] = "123456789ABCDEFGHJKLMNPQRSTUVWXYZabcdefghijkmnopqrstuvwxyz";
    constexpr size_t alphabet_size = sizeof(alphabet) - 1;
@@ -133,10 +114,6 @@ bool xmrig::WalletAddress::decode(const char *address, size_t size)
        if (memcmp(m_checksum, md, sizeof(m_checksum)) == 0) {
            m_data = { address, size };

-            if (tf_tag) {
-              m_tag = tf_tag;
-            }
-
            return true;
        }
    }
@@ -251,16 +228,6 @@ const xmrig::WalletAddress::TagInfo &xmrig::WalletAddress::tagInfo(uint64_t tag)
        { 0x54,     { Coin::GRAFT,      TESTNET,    PUBLIC,         28881,  28882 } },
        { 0x55,     { Coin::GRAFT,      TESTNET,    INTEGRATED,     28881,  28882 } },
        { 0x70,     { Coin::GRAFT,      TESTNET,    SUBADDRESS,     28881,  28882 } },
-
-        { 0x424200,     { Coin::TOWNFORGE,     MAINNET,    PUBLIC,         18881,  18882 } },
-        { 0x424201,     { Coin::TOWNFORGE,     MAINNET,    SUBADDRESS,     18881,  18882 } },
-
-        { 0x424210,     { Coin::TOWNFORGE,     TESTNET,    PUBLIC,         28881,  28882 } },
-        { 0x424211,     { Coin::TOWNFORGE,     TESTNET,    SUBADDRESS,     28881,  28882 } },
-
-        { 0x424220,     { Coin::TOWNFORGE,     STAGENET,   PUBLIC,         38881,  38882 } },
-        { 0x424221,     { Coin::TOWNFORGE,     STAGENET,   SUBADDRESS,     38881,  38882 } },
-
    };

    const auto it = tags.find(tag);
--- 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/core/config/usage.h
+++ b/src/core/config/usage.h
@@ -4,8 +4,8 @@
 * Copyright (c) 2014      Lucas Jones <https://github.com/lucasjones>
 * Copyright (c) 2014-2016 Wolf9466    <https://github.com/OhGodAPet>
 * Copyright (c) 2016      Jay D Dee   <jayddee246@gmail.com>
- * 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
@@ -64,7 +64,7 @@ static inline const std::string &usage()

 #   ifdef XMRIG_FEATURE_HTTP
    u += "      --daemon                  use daemon RPC instead of pool for solo mining\n";
-    u += "      --daemon-zmq-port=N       daemon's zmq-pub port number (only use it if daemon has it enabled)\n";
+    u += "      --daemon-zmq-port         daemon's zmq-pub port number (only use it if daemon has it enabled)\n";
    u += "      --daemon-poll-interval=N  daemon poll interval in milliseconds (default: 1000)\n";
    u += "      --daemon-job-timeout=N    daemon job timeout in milliseconds (default: 15000)\n";
    u += "      --self-select=URL         self-select block templates from URL\n";
--- a/src/crypto/cn/asm/cn_main_loop.asm
+++ b/src/crypto/cn/asm/cn_main_loop.asm
@@ -8,7 +8,6 @@ PUBLIC cnv2_mainloop_bulldozer_asm
 PUBLIC cnv2_double_mainloop_sandybridge_asm
 PUBLIC cnv2_rwz_mainloop_asm
 PUBLIC cnv2_rwz_double_mainloop_asm
-PUBLIC cnv2_upx_double_mainloop_zen3_asm

 ALIGN(64)
 cnv1_single_mainloop_asm PROC
--- a/src/crypto/cn/asm/win64/cn_main_loop.asm
+++ b/src/crypto/cn/asm/win64/cn_main_loop.asm
@@ -8,7 +8,6 @@ PUBLIC cnv2_mainloop_bulldozer_asm
 PUBLIC cnv2_double_mainloop_sandybridge_asm
 PUBLIC cnv2_rwz_mainloop_asm
 PUBLIC cnv2_rwz_double_mainloop_asm
-PUBLIC cnv2_upx_double_mainloop_zen3_asm

 ALIGN(64)
 cnv1_single_mainloop_asm PROC
--- a/src/crypto/cn/sse2neon.h
+++ b/src/crypto/cn/sse2neon.h
--- a/src/crypto/common/VirtualMemory_unix.cpp
+++ b/src/crypto/common/VirtualMemory_unix.cpp
@@ -1,7 +1,7 @@
 /* XMRig
 * Copyright (c) 2018-2020 tevador     <tevador@gmail.com>
- * Copyright (c) 2018-2023 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2023 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
@@ -17,6 +17,7 @@
 *   along with this program. If not, see <http://www.gnu.org/licenses/>.
 */

+
 #include "crypto/common/VirtualMemory.h"
 #include "backend/cpu/Cpu.h"
 #include "crypto/common/portable/mm_malloc.h"
@@ -24,7 +25,6 @@

 #include <cmath>
 #include <cstdlib>
-#include <fstream>
 #include <sys/mman.h>


@@ -84,9 +84,7 @@ static inline int hugePagesFlag(size_t size)

 bool xmrig::VirtualMemory::isHugepagesAvailable()
 {
-#   ifdef XMRIG_OS_LINUX
-    return std::ifstream("/proc/sys/vm/nr_hugepages").good() || std::ifstream("/sys/devices/system/node/node0/hugepages/hugepages-2048kB/nr_hugepages").good();
-#   elif defined(XMRIG_OS_MACOS) && defined(XMRIG_ARM)
+#   if defined(XMRIG_OS_MACOS) && defined(XMRIG_ARM)
    return false;
 #   else
    return true;
--- a/src/crypto/randomx/jit_compiler_a64.cpp
+++ b/src/crypto/randomx/jit_compiler_a64.cpp
@@ -131,8 +131,8 @@ void JitCompilerA64::generateProgram(Program& program, ProgramConfiguration& con
 	// and w16, w10, ScratchpadL3Mask64
 	emit32(0x121A0000 | 16 | (10 << 5) | ((RandomX_CurrentConfig.Log2_ScratchpadL3 - 7) << 10), code, codePos);

-	// and w17, w20, ScratchpadL3Mask64
-	emit32(0x121A0000 | 17 | (20 << 5) | ((RandomX_CurrentConfig.Log2_ScratchpadL3 - 7) << 10), code, codePos);
+	// and w17, w18, ScratchpadL3Mask64
+	emit32(0x121A0000 | 17 | (18 << 5) | ((RandomX_CurrentConfig.Log2_ScratchpadL3 - 7) << 10), code, codePos);

 	codePos = PrologueSize;
 	literalPos = ImulRcpLiteralsEnd;
@@ -148,16 +148,16 @@ void JitCompilerA64::generateProgram(Program& program, ProgramConfiguration& con
 	}

 	// Update spMix2
-	// eor w20, config.readReg2, config.readReg3
-	emit32(ARMV8A::EOR32 | 20 | (IntRegMap[config.readReg2] << 5) | (IntRegMap[config.readReg3] << 16), code, codePos);
+	// eor w18, config.readReg2, config.readReg3
+	emit32(ARMV8A::EOR32 | 18 | (IntRegMap[config.readReg2] << 5) | (IntRegMap[config.readReg3] << 16), code, codePos);

 	// Jump back to the main loop
 	const uint32_t offset = (((uint8_t*)randomx_program_aarch64_vm_instructions_end) - ((uint8_t*)randomx_program_aarch64)) - codePos;
 	emit32(ARMV8A::B | (offset / 4), code, codePos);

-	// and w20, w20, CacheLineAlignMask
+	// and w18, w18, CacheLineAlignMask
 	codePos = (((uint8_t*)randomx_program_aarch64_cacheline_align_mask1) - ((uint8_t*)randomx_program_aarch64));
-	emit32(0x121A0000 | 20 | (20 << 5) | ((RandomX_CurrentConfig.Log2_DatasetBaseSize - 7) << 10), code, codePos);
+	emit32(0x121A0000 | 18 | (18 << 5) | ((RandomX_CurrentConfig.Log2_DatasetBaseSize - 7) << 10), code, codePos);

 	// and w10, w10, CacheLineAlignMask
 	codePos = (((uint8_t*)randomx_program_aarch64_cacheline_align_mask2) - ((uint8_t*)randomx_program_aarch64));
@@ -189,8 +189,8 @@ void JitCompilerA64::generateProgramLight(Program& program, ProgramConfiguration
 	// and w16, w10, ScratchpadL3Mask64
 	emit32(0x121A0000 | 16 | (10 << 5) | ((RandomX_CurrentConfig.Log2_ScratchpadL3 - 7) << 10), code, codePos);

-	// and w17, w20, ScratchpadL3Mask64
-	emit32(0x121A0000 | 17 | (20 << 5) | ((RandomX_CurrentConfig.Log2_ScratchpadL3 - 7) << 10), code, codePos);
+	// and w17, w18, ScratchpadL3Mask64
+	emit32(0x121A0000 | 17 | (18 << 5) | ((RandomX_CurrentConfig.Log2_ScratchpadL3 - 7) << 10), code, codePos);

 	codePos = PrologueSize;
 	literalPos = ImulRcpLiteralsEnd;
@@ -206,8 +206,8 @@ void JitCompilerA64::generateProgramLight(Program& program, ProgramConfiguration
 	}

 	// Update spMix2
-	// eor w20, config.readReg2, config.readReg3
-	emit32(ARMV8A::EOR32 | 20 | (IntRegMap[config.readReg2] << 5) | (IntRegMap[config.readReg3] << 16), code, codePos);
+	// eor w18, config.readReg2, config.readReg3
+	emit32(ARMV8A::EOR32 | 18 | (IntRegMap[config.readReg2] << 5) | (IntRegMap[config.readReg3] << 16), code, codePos);

 	// Jump back to the main loop
 	const uint32_t offset = (((uint8_t*)randomx_program_aarch64_vm_instructions_end_light) - ((uint8_t*)randomx_program_aarch64)) - codePos;
@@ -477,7 +477,7 @@ void JitCompilerA64::emitAddImmediate(uint32_t dst, uint32_t src, uint32_t imm,
 	}
 	else
 	{
-		constexpr uint32_t tmp_reg = 20;
+		constexpr uint32_t tmp_reg = 18;
 		emitMovImmediate(tmp_reg, imm, code, k);

 		// add dst, src, tmp_reg
@@ -526,7 +526,7 @@ void JitCompilerA64::emitMemLoadFP(uint32_t src, Instruction& instr, uint8_t* co
 	uint32_t k = codePos;

 	uint32_t imm = instr.getImm32();
-	constexpr uint32_t tmp_reg = 19;
+	constexpr uint32_t tmp_reg = 18;

 	imm &= instr.getModMem() ? (RandomX_CurrentConfig.ScratchpadL1_Size - 1) : (RandomX_CurrentConfig.ScratchpadL2_Size - 1);
 	emitAddImmediate(tmp_reg, src, imm, code, k);
@@ -580,7 +580,7 @@ void JitCompilerA64::h_IADD_M(Instruction& instr, uint32_t& codePos)
 	const uint32_t src = IntRegMap[instr.src];
 	const uint32_t dst = IntRegMap[instr.dst];

-	constexpr uint32_t tmp_reg = 20;
+	constexpr uint32_t tmp_reg = 18;
 	emitMemLoad<tmp_reg>(dst, src, instr, code, k);

 	// add dst, dst, tmp_reg
@@ -618,7 +618,7 @@ void JitCompilerA64::h_ISUB_M(Instruction& instr, uint32_t& codePos)
 	const uint32_t src = IntRegMap[instr.src];
 	const uint32_t dst = IntRegMap[instr.dst];

-	constexpr uint32_t tmp_reg = 20;
+	constexpr uint32_t tmp_reg = 18;
 	emitMemLoad<tmp_reg>(dst, src, instr, code, k);

 	// sub dst, dst, tmp_reg
@@ -637,7 +637,7 @@ void JitCompilerA64::h_IMUL_R(Instruction& instr, uint32_t& codePos)

 	if (src == dst)
 	{
-		src = 20;
+		src = 18;
 		emitMovImmediate(src, instr.getImm32(), code, k);
 	}

@@ -655,7 +655,7 @@ void JitCompilerA64::h_IMUL_M(Instruction& instr, uint32_t& codePos)
 	const uint32_t src = IntRegMap[instr.src];
 	const uint32_t dst = IntRegMap[instr.dst];

-	constexpr uint32_t tmp_reg = 20;
+	constexpr uint32_t tmp_reg = 18;
 	emitMemLoad<tmp_reg>(dst, src, instr, code, k);

 	// sub dst, dst, tmp_reg
@@ -686,7 +686,7 @@ void JitCompilerA64::h_IMULH_M(Instruction& instr, uint32_t& codePos)
 	const uint32_t src = IntRegMap[instr.src];
 	const uint32_t dst = IntRegMap[instr.dst];

-	constexpr uint32_t tmp_reg = 20;
+	constexpr uint32_t tmp_reg = 18;
 	emitMemLoad<tmp_reg>(dst, src, instr, code, k);

 	// umulh dst, dst, tmp_reg
@@ -717,7 +717,7 @@ void JitCompilerA64::h_ISMULH_M(Instruction& instr, uint32_t& codePos)
 	const uint32_t src = IntRegMap[instr.src];
 	const uint32_t dst = IntRegMap[instr.dst];

-	constexpr uint32_t tmp_reg = 20;
+	constexpr uint32_t tmp_reg = 18;
 	emitMemLoad<tmp_reg>(dst, src, instr, code, k);

 	// smulh dst, dst, tmp_reg
@@ -735,7 +735,7 @@ void JitCompilerA64::h_IMUL_RCP(Instruction& instr, uint32_t& codePos)

 	uint32_t k = codePos;

-	constexpr uint32_t tmp_reg = 20;
+	constexpr uint32_t tmp_reg = 18;
 	const uint32_t dst = IntRegMap[instr.dst];

 	constexpr uint64_t N = 1ULL << 63;
@@ -754,9 +754,9 @@ void JitCompilerA64::h_IMUL_RCP(Instruction& instr, uint32_t& codePos)
 	literalPos -= sizeof(uint64_t);
 	*(uint64_t*)(code + literalPos) = (q << shift) + ((r << shift) / divisor);

-	if (literal_id < 12)
+	if (literal_id < 13)
 	{
-		static constexpr uint32_t literal_regs[12] = { 30 << 16, 29 << 16, 28 << 16, 27 << 16, 26 << 16, 25 << 16, 24 << 16, 23 << 16, 22 << 16, 21 << 16, 11 << 16, 0 };
+		static constexpr uint32_t literal_regs[13] = { 30 << 16, 29 << 16, 28 << 16, 27 << 16, 26 << 16, 25 << 16, 24 << 16, 23 << 16, 22 << 16, 21 << 16, 20 << 16, 11 << 16, 0 };

 		// mul dst, dst, literal_reg
 		emit32(ARMV8A::MUL | dst | (dst << 5) | literal_regs[literal_id], code, k);
@@ -794,7 +794,7 @@ void JitCompilerA64::h_IXOR_R(Instruction& instr, uint32_t& codePos)

 	if (src == dst)
 	{
-		src = 20;
+		src = 18;
 		emitMovImmediate(src, instr.getImm32(), code, k);
 	}

@@ -812,7 +812,7 @@ void JitCompilerA64::h_IXOR_M(Instruction& instr, uint32_t& codePos)
 	const uint32_t src = IntRegMap[instr.src];
 	const uint32_t dst = IntRegMap[instr.dst];

-	constexpr uint32_t tmp_reg = 20;
+	constexpr uint32_t tmp_reg = 18;
 	emitMemLoad<tmp_reg>(dst, src, instr, code, k);

 	// eor dst, dst, tmp_reg
@@ -850,7 +850,7 @@ void JitCompilerA64::h_IROL_R(Instruction& instr, uint32_t& codePos)

 	if (src != dst)
 	{
-		constexpr uint32_t tmp_reg = 20;
+		constexpr uint32_t tmp_reg = 18;

 		// sub tmp_reg, xzr, src
 		emit32(ARMV8A::SUB | tmp_reg | (31 << 5) | (src << 16), code, k);
@@ -878,7 +878,7 @@ void JitCompilerA64::h_ISWAP_R(Instruction& instr, uint32_t& codePos)

 	uint32_t k = codePos;

-	constexpr uint32_t tmp_reg = 20;
+	constexpr uint32_t tmp_reg = 18;
 	emit32(ARMV8A::MOV_REG | tmp_reg | (dst << 16), code, k);
 	emit32(ARMV8A::MOV_REG | dst | (src << 16), code, k);
 	emit32(ARMV8A::MOV_REG | src | (tmp_reg << 16), code, k);
@@ -1026,7 +1026,7 @@ void JitCompilerA64::h_CFROUND(Instruction& instr, uint32_t& codePos)

 	const uint32_t src = IntRegMap[instr.src];

-	constexpr uint32_t tmp_reg = 20;
+	constexpr uint32_t tmp_reg = 18;
 	constexpr uint32_t fpcr_tmp_reg = 8;

 	// ror tmp_reg, src, imm
@@ -1050,7 +1050,7 @@ void JitCompilerA64::h_ISTORE(Instruction& instr, uint32_t& codePos)

 	const uint32_t src = IntRegMap[instr.src];
 	const uint32_t dst = IntRegMap[instr.dst];
-	constexpr uint32_t tmp_reg = 20;
+	constexpr uint32_t tmp_reg = 18;

 	uint32_t imm = instr.getImm32();

--- a/src/crypto/randomx/jit_compiler_a64_static.S
+++ b/src/crypto/randomx/jit_compiler_a64_static.S
@@ -72,9 +72,9 @@
 # x15 -> "r7"
 # x16 -> spAddr0
 # x17 -> spAddr1
-# x18 -> unused (platform register, don't touch it)
+# x18 -> temporary
 # x19 -> temporary
-# x20 -> temporary
+# x20 -> literal for IMUL_RCP
 # x21 -> literal for IMUL_RCP
 # x22 -> literal for IMUL_RCP
 # x23 -> literal for IMUL_RCP
@@ -109,7 +109,7 @@ DECL(randomx_program_aarch64):
 	# Save callee-saved registers
 	sub	sp, sp, 192
 	stp	x16, x17, [sp]
-	str	x19, [sp, 16]
+	stp	x18, x19, [sp, 16]
 	stp	x20, x21, [sp, 32]
 	stp	x22, x23, [sp, 48]
 	stp	x24, x25, [sp, 64]
@@ -164,6 +164,7 @@ DECL(randomx_program_aarch64):
 	# Read literals
 	ldr	x0, literal_x0
 	ldr	x11, literal_x11
+	ldr	x20, literal_x20
 	ldr	x21, literal_x21
 	ldr	x22, literal_x22
 	ldr	x23, literal_x23
@@ -195,11 +196,11 @@ DECL(randomx_program_aarch64):
 DECL(randomx_program_aarch64_main_loop):
 	# spAddr0 = spMix1 & ScratchpadL3Mask64;
 	# spAddr1 = (spMix1 >> 32) & ScratchpadL3Mask64;
-	lsr	x20, x10, 32
+	lsr	x18, x10, 32

 	# Actual mask will be inserted by JIT compiler
 	and	w16, w10, 1
-	and	w17, w20, 1
+	and	w17, w18, 1

 	# x16 = scratchpad + spAddr0
 	# x17 = scratchpad + spAddr1
@@ -207,31 +208,31 @@ DECL(randomx_program_aarch64_main_loop):
 	add	x17, x17, x2

 	# xor integer registers with scratchpad data (spAddr0)
-	ldp	x20, x19, [x16]
-	eor	x4, x4, x20
+	ldp	x18, x19, [x16]
+	eor	x4, x4, x18
 	eor	x5, x5, x19
-	ldp	x20, x19, [x16, 16]
-	eor	x6, x6, x20
+	ldp	x18, x19, [x16, 16]
+	eor	x6, x6, x18
 	eor	x7, x7, x19
-	ldp	x20, x19, [x16, 32]
-	eor	x12, x12, x20
+	ldp	x18, x19, [x16, 32]
+	eor	x12, x12, x18
 	eor	x13, x13, x19
-	ldp	x20, x19, [x16, 48]
-	eor	x14, x14, x20
+	ldp	x18, x19, [x16, 48]
+	eor	x14, x14, x18
 	eor	x15, x15, x19

 	# Load group F registers (spAddr1)
-	ldpsw	x20, x19, [x17]
-	ins	v16.d[0], x20
+	ldpsw	x18, x19, [x17]
+	ins	v16.d[0], x18
 	ins	v16.d[1], x19
-	ldpsw	x20, x19, [x17, 8]
-	ins	v17.d[0], x20
+	ldpsw	x18, x19, [x17, 8]
+	ins	v17.d[0], x18
 	ins	v17.d[1], x19
-	ldpsw	x20, x19, [x17, 16]
-	ins	v18.d[0], x20
+	ldpsw	x18, x19, [x17, 16]
+	ins	v18.d[0], x18
 	ins	v18.d[1], x19
-	ldpsw	x20, x19, [x17, 24]
-	ins	v19.d[0], x20
+	ldpsw	x18, x19, [x17, 24]
+	ins	v19.d[0], x18
 	ins	v19.d[1], x19
 	scvtf	v16.2d, v16.2d
 	scvtf	v17.2d, v17.2d
@@ -239,17 +240,17 @@ DECL(randomx_program_aarch64_main_loop):
 	scvtf	v19.2d, v19.2d

 	# Load group E registers (spAddr1)
-	ldpsw	x20, x19, [x17, 32]
-	ins	v20.d[0], x20
+	ldpsw	x18, x19, [x17, 32]
+	ins	v20.d[0], x18
 	ins	v20.d[1], x19
-	ldpsw	x20, x19, [x17, 40]
-	ins	v21.d[0], x20
+	ldpsw	x18, x19, [x17, 40]
+	ins	v21.d[0], x18
 	ins	v21.d[1], x19
-	ldpsw	x20, x19, [x17, 48]
-	ins	v22.d[0], x20
+	ldpsw	x18, x19, [x17, 48]
+	ins	v22.d[0], x18
 	ins	v22.d[1], x19
-	ldpsw	x20, x19, [x17, 56]
-	ins	v23.d[0], x20
+	ldpsw	x18, x19, [x17, 56]
+	ins	v23.d[0], x18
 	ins	v23.d[1], x19
 	scvtf	v20.2d, v20.2d
 	scvtf	v21.2d, v21.2d
@@ -272,6 +273,7 @@ DECL(randomx_program_aarch64_vm_instructions):

 literal_x0:  .fill 1,8,0
 literal_x11: .fill 1,8,0
+literal_x20: .fill 1,8,0
 literal_x21: .fill 1,8,0
 literal_x22: .fill 1,8,0
 literal_x23: .fill 1,8,0
@@ -307,17 +309,17 @@ DECL(randomx_program_aarch64_vm_instructions_end):
 	lsr	x10, x9, 32

 	# mx ^= r[readReg2] ^ r[readReg3];
-	eor	x9, x9, x20
+	eor	x9, x9, x18

 	# Calculate dataset pointer for dataset prefetch
-	mov	w20, w9
+	mov	w18, w9
 DECL(randomx_program_aarch64_cacheline_align_mask1):
 	# Actual mask will be inserted by JIT compiler
-	and	x20, x20, 1
-	add	x20, x20, x1
+	and	x18, x18, 1
+	add	x18, x18, x1

 	# Prefetch dataset data
-	prfm	pldl2strm, [x20]
+	prfm	pldl2strm, [x18]

 	# mx <-> ma
 	ror	x9, x9, 32
@@ -329,17 +331,17 @@ DECL(randomx_program_aarch64_cacheline_align_mask2):

 DECL(randomx_program_aarch64_xor_with_dataset_line):
 	# xor integer registers with dataset data
-	ldp	x20, x19, [x10]
-	eor	x4, x4, x20
+	ldp	x18, x19, [x10]
+	eor	x4, x4, x18
 	eor	x5, x5, x19
-	ldp	x20, x19, [x10, 16]
-	eor	x6, x6, x20
+	ldp	x18, x19, [x10, 16]
+	eor	x6, x6, x18
 	eor	x7, x7, x19
-	ldp	x20, x19, [x10, 32]
-	eor	x12, x12, x20
+	ldp	x18, x19, [x10, 32]
+	eor	x12, x12, x18
 	eor	x13, x13, x19
-	ldp	x20, x19, [x10, 48]
-	eor	x14, x14, x20
+	ldp	x18, x19, [x10, 48]
+	eor	x14, x14, x18
 	eor	x15, x15, x19

 DECL(randomx_program_aarch64_update_spMix1):
@@ -382,7 +384,7 @@ DECL(randomx_program_aarch64_update_spMix1):

 	# Restore callee-saved registers
 	ldp	x16, x17, [sp]
-	ldr	x19, [sp, 16]
+	ldp	x18, x19, [sp, 16]
 	ldp	x20, x21, [sp, 32]
 	ldp	x22, x23, [sp, 48]
 	ldp	x24, x25, [sp, 64]
@@ -403,7 +405,7 @@ DECL(randomx_program_aarch64_vm_instructions_end_light):
 	stp	x2, x30, [sp, 80]

 	# mx ^= r[readReg2] ^ r[readReg3];
-	eor	x9, x9, x20
+	eor	x9, x9, x18

 	# mx <-> ma
 	ror	x9, x9, 32
@@ -445,8 +447,8 @@ DECL(randomx_program_aarch64_light_dataset_offset):
 # x3 -> end item

 DECL(randomx_init_dataset_aarch64):
-	# Save x20 (used as temporary, but must be saved to not break ABI) and x30 (return address)
-	stp	x20, x30, [sp, -16]!
+	# Save x30 (return address)
+	str	x30, [sp, -16]!

 	# Load pointer to cache memory
 	ldr	x0, [x0]
@@ -458,8 +460,8 @@ DECL(randomx_init_dataset_aarch64_main_loop):
 	cmp	x2, x3
 	bne	DECL(randomx_init_dataset_aarch64_main_loop)

-	# Restore x20 and x30
-	ldp	x20, x30, [sp], 16
+	# Restore x30 (return address)
+	ldr	x30, [sp], 16

 	ret

--- a/src/version.h
+++ b/src/version.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-2023 SChernykh   <https://github.com/SChernykh>
+ * Copyright (c) 2016-2023 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,14 +22,14 @@
 #define APP_ID        "xmrig"
 #define APP_NAME      "XMRig"
 #define APP_DESC      "XMRig miner"
-#define APP_VERSION   "6.21.1"
+#define APP_VERSION   "6.20.1-dev"
 #define APP_DOMAIN    "xmrig.com"
 #define APP_SITE      "www.xmrig.com"
-#define APP_COPYRIGHT "Copyright (C) 2016-2024 xmrig.com"
+#define APP_COPYRIGHT "Copyright (C) 2016-2023 xmrig.com"
 #define APP_KIND      "miner"

 #define APP_VER_MAJOR  6
-#define APP_VER_MINOR  21
+#define APP_VER_MINOR  20
 #define APP_VER_PATCH  1

 #ifdef _MSC_VER
@@ -52,39 +52,4 @@
 #   endif
 #endif

-#ifdef XMRIG_OS_WIN
-#    define APP_OS "Windows"
-#elif defined XMRIG_OS_IOS
-#    define APP_OS "iOS"
-#elif defined XMRIG_OS_MACOS
-#    define APP_OS "macOS"
-#elif defined XMRIG_OS_ANDROID
-#    define APP_OS "Android"
-#elif defined XMRIG_OS_LINUX
-#    define APP_OS "Linux"
-#elif defined XMRIG_OS_FREEBSD
-#    define APP_OS "FreeBSD"
-#else
-#    define APP_OS "Unknown OS"
-#endif
-
-#define STR(X) #X
-#define STR2(X) STR(X)
-
-#ifdef XMRIG_ARM
-#   define APP_ARCH "ARMv" STR2(XMRIG_ARM)
-#else
-#   if defined(__x86_64__) || defined(__amd64__) || defined(_M_X64) || defined(_M_AMD64)
-#       define APP_ARCH "x86-64"
-#   else
-#       define APP_ARCH "x86"
-#   endif
-#endif
-
-#ifdef XMRIG_64_BIT
-#   define APP_BITS "64 bit"
-#else
-#   define APP_BITS "32 bit"
-#endif
-
 #endif // XMRIG_VERSION_H
Author	SHA1	Message	Date
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