1
0
mirror of https://github.com/xmrig/xmrig.git synced 2025-12-12 17:52:48 -05:00

Compare commits

..

3 Commits

Author SHA1 Message Date
Tony Butler
a11fa20268 Merge 753e63cd96 into e855723cd9 2023-09-01 13:28:19 -07: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
34 changed files with 3594 additions and 3775 deletions

View File

@@ -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 # v6.20.0
- Added new ARM CPU names. - Added new ARM CPU names.
- [#2394](https://github.com/xmrig/xmrig/pull/2394) Added new CMake options `ARM_V8` and `ARM_V7`. - [#2394](https://github.com/xmrig/xmrig/pull/2394) Added new CMake options `ARM_V8` and `ARM_V7`.

View File

@@ -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/ Offical HTTP client for API: http://workers.xmrig.info/
Example configuration: Example configuration, used in Curl examples below:
```json ```json
"api": { "api": {
@@ -12,11 +12,11 @@ Example configuration:
"worker-id": null, "worker-id": null,
}, },
"http": { "http": {
"enabled": false, "enabled": true,
"host": "127.0.0.1", "host": "127.0.0.1",
"port": 0, "port": 44444,
"access-token": null, "access-token": "SECRET",
"restricted": true "restricted": false
} }
``` ```
@@ -37,30 +37,78 @@ Versions before 2.15 was use another options for API https://github.com/xmrig/xm
## Endpoints ## 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 ## 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`). 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. Update current miner configuration. Common use case, get current configuration, make changes, and upload it to miner.
Curl example: 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`

78
doc/api/2/backends.json Normal file
View File

@@ -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]
}
}
]
}
]

136
doc/api/2/config.json Normal file
View File

@@ -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
}

76
doc/api/2/summary.json Normal file
View File

@@ -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]
}

View File

@@ -138,93 +138,6 @@ __kernel void blake2b_initial_hash(__global void *out, __global const void* bloc
t[7] = hash[7]; 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 in_len 256
#define out_len 32 #define out_len 32

File diff suppressed because it is too large Load Diff

View File

@@ -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);
}

View File

@@ -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 */

View File

@@ -80,7 +80,6 @@ if (WITH_OPENCL)
if (WITH_RANDOMX) if (WITH_RANDOMX)
list(APPEND HEADERS_BACKEND_OPENCL list(APPEND HEADERS_BACKEND_OPENCL
src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h 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/Blake2bInitialHashKernel.h
src/backend/opencl/kernels/rx/ExecuteVmKernel.h src/backend/opencl/kernels/rx/ExecuteVmKernel.h
src/backend/opencl/kernels/rx/FillAesKernel.h src/backend/opencl/kernels/rx/FillAesKernel.h
@@ -97,7 +96,6 @@ if (WITH_OPENCL)
list(APPEND SOURCES_BACKEND_OPENCL list(APPEND SOURCES_BACKEND_OPENCL
src/backend/opencl/generators/ocl_generic_rx_generator.cpp src/backend/opencl/generators/ocl_generic_rx_generator.cpp
src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.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/Blake2bInitialHashKernel.cpp
src/backend/opencl/kernels/rx/ExecuteVmKernel.cpp src/backend/opencl/kernels/rx/ExecuteVmKernel.cpp
src/backend/opencl/kernels/rx/FillAesKernel.cpp src/backend/opencl/kernels/rx/FillAesKernel.cpp

View File

@@ -25,7 +25,6 @@
#include "backend/opencl/runners/OclRxBaseRunner.h" #include "backend/opencl/runners/OclRxBaseRunner.h"
#include "backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h" #include "backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h"
#include "backend/opencl/kernels/rx/Blake2bInitialHashKernel.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/FillAesKernel.h"
#include "backend/opencl/kernels/rx/FindSharesKernel.h" #include "backend/opencl/kernels/rx/FindSharesKernel.h"
#include "backend/opencl/kernels/rx/HashAesKernel.h" #include "backend/opencl/kernels/rx/HashAesKernel.h"
@@ -72,7 +71,6 @@ xmrig::OclRxBaseRunner::~OclRxBaseRunner()
delete m_fillAes4Rx4_entropy; delete m_fillAes4Rx4_entropy;
delete m_hashAes1Rx4; delete m_hashAes1Rx4;
delete m_blake2b_initial_hash; delete m_blake2b_initial_hash;
delete m_blake2b_initial_hash_double;
delete m_blake2b_hash_registers_32; delete m_blake2b_hash_registers_32;
delete m_blake2b_hash_registers_64; delete m_blake2b_hash_registers_64;
delete m_find_shares; delete m_find_shares;
@@ -89,28 +87,12 @@ void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
{ {
static const uint32_t zero = 0; static const uint32_t zero = 0;
if (m_jobSize <= 128) {
m_blake2b_initial_hash->setNonce(nonce); m_blake2b_initial_hash->setNonce(nonce);
}
else if (m_jobSize <= 256) {
m_blake2b_initial_hash_double->setNonce(nonce);
}
else {
hashOutput[0xFF] = 0;
return;
}
m_find_shares->setNonce(nonce); m_find_shares->setNonce(nonce);
enqueueWriteBuffer(m_output, CL_FALSE, sizeof(cl_uint) * 0xFF, sizeof(uint32_t), &zero); 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); m_blake2b_initial_hash->enqueue(m_queue, m_intensity);
}
else {
m_blake2b_initial_hash_double->enqueue(m_queue, m_intensity);
}
m_fillAes1Rx4_scratchpad->enqueue(m_queue, m_intensity); m_fillAes1Rx4_scratchpad->enqueue(m_queue, m_intensity);
const uint32_t programCount = RxAlgo::programCount(m_algorithm); 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); enqueueWriteBuffer(m_input, CL_TRUE, 0, Job::kMaxBlobSize, blob);
m_jobSize = job.size();
m_blake2b_initial_hash->setBlobSize(job.size()); m_blake2b_initial_hash->setBlobSize(job.size());
m_blake2b_initial_hash_double->setBlobSize(job.size());
m_find_shares->setTarget(job.target()); 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 = new Blake2bInitialHashKernel(m_program);
m_blake2b_initial_hash->setArgs(m_hashes, m_input); 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_32 = new Blake2bHashRegistersKernel(m_program, "blake2b_hash_registers_32");
m_blake2b_hash_registers_64 = new Blake2bHashRegistersKernel(m_program, "blake2b_hash_registers_64"); m_blake2b_hash_registers_64 = new Blake2bHashRegistersKernel(m_program, "blake2b_hash_registers_64");

View File

@@ -35,7 +35,6 @@ namespace xmrig {
class Blake2bHashRegistersKernel; class Blake2bHashRegistersKernel;
class Blake2bInitialHashKernel; class Blake2bInitialHashKernel;
class Blake2bInitialHashDoubleKernel;
class FillAesKernel; class FillAesKernel;
class FindSharesKernel; class FindSharesKernel;
class HashAesKernel; class HashAesKernel;
@@ -62,7 +61,6 @@ protected:
Blake2bHashRegistersKernel *m_blake2b_hash_registers_32 = nullptr; Blake2bHashRegistersKernel *m_blake2b_hash_registers_32 = nullptr;
Blake2bHashRegistersKernel *m_blake2b_hash_registers_64 = nullptr; Blake2bHashRegistersKernel *m_blake2b_hash_registers_64 = nullptr;
Blake2bInitialHashKernel *m_blake2b_initial_hash = nullptr; Blake2bInitialHashKernel *m_blake2b_initial_hash = nullptr;
Blake2bInitialHashDoubleKernel *m_blake2b_initial_hash_double = nullptr;
Buffer m_seed; Buffer m_seed;
cl_mem m_dataset = nullptr; cl_mem m_dataset = nullptr;
cl_mem m_entropy = nullptr; cl_mem m_entropy = nullptr;
@@ -75,8 +73,6 @@ protected:
HashAesKernel *m_hashAes1Rx4 = nullptr; HashAesKernel *m_hashAes1Rx4 = nullptr;
uint32_t m_gcn_version = 12; uint32_t m_gcn_version = 12;
uint32_t m_worksize = 8; uint32_t m_worksize = 8;
size_t m_jobSize = 0;
}; };

View File

@@ -39,7 +39,6 @@
#include <thread> #include <thread>
#include <iostream>
namespace xmrig { namespace xmrig {
@@ -81,8 +80,7 @@ static rapidjson::Value getResources(rapidjson::Document &doc)
xmrig::Api::Api(Base *base) : xmrig::Api::Api(Base *base) :
m_base(base), m_base(base),
m_timestamp(Chrono::currentMSecsSinceEpoch()), m_timestamp(Chrono::currentMSecsSinceEpoch())
m_httpd(nullptr)
{ {
base->addListener(this); base->addListener(this);
@@ -93,11 +91,7 @@ xmrig::Api::Api(Base *base) :
xmrig::Api::~Api() xmrig::Api::~Api()
{ {
# ifdef XMRIG_FEATURE_HTTP # ifdef XMRIG_FEATURE_HTTP
if (m_httpd) {
m_httpd->stop();
delete m_httpd; delete m_httpd;
m_httpd = nullptr; // Ensure the pointer is set to nullptr after deletion
}
# endif # endif
} }
@@ -115,14 +109,8 @@ void xmrig::Api::start()
genWorkerId(m_base->config()->apiWorkerId()); genWorkerId(m_base->config()->apiWorkerId());
# ifdef XMRIG_FEATURE_HTTP # ifdef XMRIG_FEATURE_HTTP
if (!m_httpd) {
m_httpd = new Httpd(m_base); m_httpd = new Httpd(m_base);
if (!m_httpd->start()) { m_httpd->start();
std::cerr << "HTTP server failed to start." << std::endl;
delete m_httpd; // Properly handle failure to start
m_httpd = nullptr;
}
}
# endif # endif
} }
@@ -130,9 +118,7 @@ void xmrig::Api::start()
void xmrig::Api::stop() void xmrig::Api::stop()
{ {
# ifdef XMRIG_FEATURE_HTTP # ifdef XMRIG_FEATURE_HTTP
if (m_httpd) {
m_httpd->stop(); m_httpd->stop();
}
# endif # endif
} }
@@ -140,16 +126,14 @@ void xmrig::Api::stop()
void xmrig::Api::tick() void xmrig::Api::tick()
{ {
# ifdef XMRIG_FEATURE_HTTP # 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; return;
} }
if (++m_ticks % 10 == 0) { if (++m_ticks % 10 == 0) {
m_ticks = 0; m_ticks = 0;
if (m_httpd) {
m_httpd->start(); m_httpd->start();
} }
}
# endif # endif
} }

View File

@@ -52,6 +52,8 @@ public:
enum RequestType { enum RequestType {
REQ_UNKNOWN, REQ_UNKNOWN,
REQ_SUMMARY, REQ_SUMMARY,
REQ_BACKENDS,
REQ_CONFIG,
REQ_JSON_RPC REQ_JSON_RPC
}; };

View File

@@ -60,7 +60,7 @@ protected:
inline Source source() const override { return m_source; } inline Source source() const override { return m_source; }
inline void done(int) override { m_state = STATE_DONE; } inline void done(int) override { m_state = STATE_DONE; }
int m_version = 1; int m_version = 0;
RequestType m_type = REQ_UNKNOWN; RequestType m_type = REQ_UNKNOWN;
State m_state = STATE_NEW; State m_state = STATE_NEW;
String m_rpcMethod; String m_rpcMethod;

View File

@@ -67,10 +67,33 @@ xmrig::HttpApiRequest::HttpApiRequest(const HttpData &req, bool restricted) :
m_res(req.id()), m_res(req.id()),
m_url(req.url.c_str()) m_url(req.url.c_str())
{ {
if (method() == METHOD_GET) { if (url().size() > 4 && memcmp(url().data(), "/", 1) == 0 && memcmp(url().data()+2, "/", 1) == 0) {
if (url() == "/1/summary" || url() == "/2/summary" || url() == "/api.json") { 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; 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") { if (method() == METHOD_POST && url() == "/json_rpc") {
@@ -94,12 +117,6 @@ xmrig::HttpApiRequest::HttpApiRequest(const HttpData &req, bool restricted) :
return; return;
} }
if (url().size() > 4) {
if (memcmp(url().data(), "/2/", 3) == 0) {
m_version = 2;
}
}
} }

View File

@@ -54,7 +54,6 @@ static const CoinInfo coinInfo[] = {
{ Algorithm::KAWPOW_RVN, "RVN", "Ravencoin", 0, 0, BLUE_BG_BOLD( WHITE_BOLD_S " raven ") }, { 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_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, "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 ") },
}; };

View File

@@ -40,7 +40,6 @@ public:
RAVEN, RAVEN,
WOWNERO, WOWNERO,
ZEPHYR, ZEPHYR,
TOWNFORGE,
MAX MAX
}; };

View File

@@ -45,13 +45,6 @@
#ifdef XMRIG_FEATURE_API #ifdef XMRIG_FEATURE_API
# include "base/api/Api.h" # include "base/api/Api.h"
# include "base/api/interfaces/IApiRequest.h" # include "base/api/interfaces/IApiRequest.h"
namespace xmrig {
static const char *kConfigPathV1 = "/1/config";
static const char *kConfigPathV2 = "/2/config";
} // namespace xmrig
#endif #endif
@@ -317,7 +310,7 @@ void xmrig::Base::onFileChanged(const String &fileName)
void xmrig::Base::onRequest(IApiRequest &request) void xmrig::Base::onRequest(IApiRequest &request)
{ {
if (request.method() == IApiRequest::METHOD_GET) { if (request.method() == IApiRequest::METHOD_GET) {
if (request.url() == kConfigPathV1 || request.url() == kConfigPathV2) { if (request.type() == IApiRequest::REQ_CONFIG) {
if (request.isRestricted()) { if (request.isRestricted()) {
return request.done(403); 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) { 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(); request.accept();
if (!reload(request.json())) { if (!reload(request.json())) {

View File

@@ -589,7 +589,7 @@ void xmrig::Client::handshake()
if (isTLS()) { if (isTLS()) {
m_expire = Chrono::steadyMSecs() + kResponseTimeout; m_expire = Chrono::steadyMSecs() + kResponseTimeout;
m_tls->handshake(m_pool.isSNI() ? m_pool.host().data() : nullptr); m_tls->handshake();
} }
else else
# endif # endif

View File

@@ -77,7 +77,6 @@ const char *Pool::kSelfSelect = "self-select";
const char *Pool::kSOCKS5 = "socks5"; const char *Pool::kSOCKS5 = "socks5";
const char *Pool::kSubmitToOrigin = "submit-to-origin"; const char *Pool::kSubmitToOrigin = "submit-to-origin";
const char *Pool::kTls = "tls"; const char *Pool::kTls = "tls";
const char *Pool::kSni = "sni";
const char *Pool::kUrl = "url"; const char *Pool::kUrl = "url";
const char *Pool::kUser = "user"; const char *Pool::kUser = "user";
const char *Pool::kSpendSecretKey = "spend-secret-key"; 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_ENABLED, Json::getBool(object, kEnabled, true));
m_flags.set(FLAG_NICEHASH, Json::getBool(object, kNicehash) || m_url.host().contains(kNicehashHost)); 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_TLS, Json::getBool(object, kTls) || m_url.isTLS());
m_flags.set(FLAG_SNI, Json::getBool(object, kSni));
setKeepAlive(Json::getValue(object, kKeepalive)); 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(kEnabled), m_flags.test(FLAG_ENABLED), allocator);
obj.AddMember(StringRef(kTls), isTLS(), allocator); obj.AddMember(StringRef(kTls), isTLS(), allocator);
obj.AddMember(StringRef(kSni), isSNI(), allocator);
obj.AddMember(StringRef(kFingerprint), m_fingerprint.toJSON(), allocator); obj.AddMember(StringRef(kFingerprint), m_fingerprint.toJSON(), allocator);
obj.AddMember(StringRef(kDaemon), m_mode == MODE_DAEMON, allocator); obj.AddMember(StringRef(kDaemon), m_mode == MODE_DAEMON, allocator);
obj.AddMember(StringRef(kSOCKS5), m_proxy.toJSON(doc), allocator); obj.AddMember(StringRef(kSOCKS5), m_proxy.toJSON(doc), allocator);

View File

@@ -1,7 +1,7 @@
/* XMRig /* XMRig
* Copyright (c) 2019 Howard Chu <https://github.com/hyc> * Copyright (c) 2019 Howard Chu <https://github.com/hyc>
* Copyright (c) 2018-2024 SChernykh <https://github.com/SChernykh> * Copyright (c) 2018-2021 SChernykh <https://github.com/SChernykh>
* Copyright (c) 2016-2024 XMRig <https://github.com/xmrig>, <support@xmrig.com> * Copyright (c) 2016-2021 XMRig <https://github.com/xmrig>, <support@xmrig.com>
* *
* This program is free software: you can redistribute it and/or modify * 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 * it under the terms of the GNU General Public License as published by
@@ -70,7 +70,6 @@ public:
static const char *kSOCKS5; static const char *kSOCKS5;
static const char *kSubmitToOrigin; static const char *kSubmitToOrigin;
static const char *kTls; static const char *kTls;
static const char *kSni;
static const char *kUrl; static const char *kUrl;
static const char *kUser; static const char *kUser;
static const char* kSpendSecretKey; static const char* kSpendSecretKey;
@@ -96,7 +95,6 @@ public:
inline bool isNicehash() const { return m_flags.test(FLAG_NICEHASH); } 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 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 bool isValid() const { return m_url.isValid(); }
inline const Algorithm &algorithm() const { return m_algorithm; } inline const Algorithm &algorithm() const { return m_algorithm; }
inline const Coin &coin() const { return m_coin; } inline const Coin &coin() const { return m_coin; }
@@ -140,7 +138,6 @@ private:
FLAG_ENABLED, FLAG_ENABLED,
FLAG_NICEHASH, FLAG_NICEHASH,
FLAG_TLS, FLAG_TLS,
FLAG_SNI,
FLAG_MAX FLAG_MAX
}; };

View File

@@ -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); m_ssl = SSL_new(m_ctx);
assert(m_ssl != nullptr); assert(m_ssl != nullptr);
@@ -69,10 +69,6 @@ bool xmrig::Client::Tls::handshake(const char* servername)
return false; return false;
} }
if (servername) {
SSL_set_tlsext_host_name(m_ssl, servername);
}
SSL_set_connect_state(m_ssl); SSL_set_connect_state(m_ssl);
SSL_set_bio(m_ssl, m_read, m_write); SSL_set_bio(m_ssl, m_read, m_write);
SSL_do_handshake(m_ssl); SSL_do_handshake(m_ssl);

View File

@@ -42,7 +42,7 @@ public:
Tls(Client *client); Tls(Client *client);
~Tls(); ~Tls();
bool handshake(const char* servername); bool handshake();
bool send(const char *data, size_t size); bool send(const char *data, size_t size);
const char *fingerprint() const; const char *fingerprint() const;
const char *version() const; const char *version() const;

View File

@@ -1,8 +1,8 @@
/* XMRig /* XMRig
* Copyright (c) 2012-2013 The Cryptonote developers * Copyright (c) 2012-2013 The Cryptonote developers
* Copyright (c) 2014-2021 The Monero Project * Copyright (c) 2014-2021 The Monero Project
* Copyright (c) 2018-2023 SChernykh <https://github.com/SChernykh> * Copyright (c) 2018-2021 SChernykh <https://github.com/SChernykh>
* Copyright (c) 2016-2023 XMRig <https://github.com/xmrig>, <support@xmrig.com> * Copyright (c) 2016-2021 XMRig <https://github.com/xmrig>, <support@xmrig.com>
* *
* This program is free software: you can redistribute it and/or modify * 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 * 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) { if (m_coin == Coin::ZEPHYR) {
uint8_t pricing_record[120]; uint8_t pricing_record[24];
ar(pricing_record); ar(pricing_record);
} }
@@ -207,11 +207,7 @@ bool xmrig::BlockTemplate::parse(bool hashes)
setOffset(MINER_TX_PREFIX_OFFSET, ar.index()); setOffset(MINER_TX_PREFIX_OFFSET, ar.index());
ar(m_txVersion); ar(m_txVersion);
if (m_coin != Coin::TOWNFORGE) {
ar(m_unlockTime); ar(m_unlockTime);
}
ar(m_numInputs); ar(m_numInputs);
// must be 1 input // must be 1 input
@@ -229,12 +225,8 @@ bool xmrig::BlockTemplate::parse(bool hashes)
ar(m_height); ar(m_height);
ar(m_numOutputs); ar(m_numOutputs);
if (m_coin == Coin::ZEPHYR) { const uint64_t expected_outputs = (m_coin == Coin::ZEPHYR) ? 2 : 1;
if (m_numOutputs < 2) { if (m_numOutputs != expected_outputs) {
return false;
}
}
else if (m_numOutputs != 1) {
return false; return false;
} }
@@ -260,7 +252,6 @@ bool xmrig::BlockTemplate::parse(bool hashes)
ar.skip(asset_type_len); ar.skip(asset_type_len);
ar(m_viewTag); ar(m_viewTag);
for (uint64_t k = 1; k < m_numOutputs; ++k) {
uint64_t amount2; uint64_t amount2;
ar(amount2); ar(amount2);
@@ -279,15 +270,10 @@ bool xmrig::BlockTemplate::parse(bool hashes)
uint8_t view_tag2; uint8_t view_tag2;
ar(view_tag2); ar(view_tag2);
} }
}
else if (m_outputType == 3) { else if (m_outputType == 3) {
ar(m_viewTag); ar(m_viewTag);
} }
if (m_coin == Coin::TOWNFORGE) {
ar(m_unlockTime);
}
ar(m_extraSize); ar(m_extraSize);
setOffset(TX_EXTRA_OFFSET, ar.index()); setOffset(TX_EXTRA_OFFSET, ar.index());
@@ -343,11 +329,6 @@ bool xmrig::BlockTemplate::parse(bool hashes)
uint8_t vin_rct_type = 0; uint8_t vin_rct_type = 0;
ar(vin_rct_type); 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) // must be RCTTypeNull (0)
if (vin_rct_type != 0) { if (vin_rct_type != 0) {
return false; return false;

View File

@@ -1,8 +1,8 @@
/* XMRig /* XMRig
* Copyright (c) 2012-2013 The Cryptonote developers * Copyright (c) 2012-2013 The Cryptonote developers
* Copyright (c) 2014-2021 The Monero Project * Copyright (c) 2014-2021 The Monero Project
* Copyright (c) 2018-2023 SChernykh <https://github.com/SChernykh> * Copyright (c) 2018-2021 SChernykh <https://github.com/SChernykh>
* Copyright (c) 2016-2023 XMRig <https://github.com/xmrig>, <support@xmrig.com> * Copyright (c) 2016-2021 XMRig <https://github.com/xmrig>, <support@xmrig.com>
* *
* This program is free software: you can redistribute it and/or modify * 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 * 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) 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 std::array<int, 9> block_sizes{ 0, 2, 3, 5, 6, 7, 9, 10, 11 };
static constexpr char alphabet[] = "123456789ABCDEFGHJKLMNPQRSTUVWXYZabcdefghijkmnopqrstuvwxyz"; static constexpr char alphabet[] = "123456789ABCDEFGHJKLMNPQRSTUVWXYZabcdefghijkmnopqrstuvwxyz";
constexpr size_t alphabet_size = sizeof(alphabet) - 1; 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) { if (memcmp(m_checksum, md, sizeof(m_checksum)) == 0) {
m_data = { address, size }; m_data = { address, size };
if (tf_tag) {
m_tag = tf_tag;
}
return true; return true;
} }
} }
@@ -251,16 +228,6 @@ const xmrig::WalletAddress::TagInfo &xmrig::WalletAddress::tagInfo(uint64_t tag)
{ 0x54, { Coin::GRAFT, TESTNET, PUBLIC, 28881, 28882 } }, { 0x54, { Coin::GRAFT, TESTNET, PUBLIC, 28881, 28882 } },
{ 0x55, { Coin::GRAFT, TESTNET, INTEGRATED, 28881, 28882 } }, { 0x55, { Coin::GRAFT, TESTNET, INTEGRATED, 28881, 28882 } },
{ 0x70, { Coin::GRAFT, TESTNET, SUBADDRESS, 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); const auto it = tags.find(tag);

View File

@@ -689,7 +689,7 @@ void xmrig::Miner::onRequest(IApiRequest &request)
d_ptr->getMiner(request.reply(), request.doc(), request.version()); d_ptr->getMiner(request.reply(), request.doc(), request.version());
d_ptr->getHashrate(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(); request.accept();
d_ptr->getBackends(request.reply(), request.doc()); d_ptr->getBackends(request.reply(), request.doc());
@@ -711,6 +711,12 @@ void xmrig::Miner::onRequest(IApiRequest &request)
stop(); stop();
} }
else if (request.rpcMethod() == "start") {
request.accept();
const auto config = d_ptr->controller->config();
onConfigChanged(config, config);
}
} }
for (IBackend *backend : d_ptr->backends) { for (IBackend *backend : d_ptr->backends) {

View File

@@ -4,8 +4,8 @@
* Copyright (c) 2014 Lucas Jones <https://github.com/lucasjones> * Copyright (c) 2014 Lucas Jones <https://github.com/lucasjones>
* Copyright (c) 2014-2016 Wolf9466 <https://github.com/OhGodAPet> * Copyright (c) 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
* Copyright (c) 2016 Jay D Dee <jayddee246@gmail.com> * Copyright (c) 2016 Jay D Dee <jayddee246@gmail.com>
* Copyright (c) 2018-2024 SChernykh <https://github.com/SChernykh> * Copyright (c) 2018-2021 SChernykh <https://github.com/SChernykh>
* Copyright (c) 2016-2024 XMRig <https://github.com/xmrig>, <support@xmrig.com> * Copyright (c) 2016-2021 XMRig <https://github.com/xmrig>, <support@xmrig.com>
* *
* This program is free software: you can redistribute it and/or modify * 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 * 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 # ifdef XMRIG_FEATURE_HTTP
u += " --daemon use daemon RPC instead of pool for solo mining\n"; 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-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 += " --daemon-job-timeout=N daemon job timeout in milliseconds (default: 15000)\n";
u += " --self-select=URL self-select block templates from URL\n"; u += " --self-select=URL self-select block templates from URL\n";

View File

@@ -8,7 +8,6 @@ PUBLIC cnv2_mainloop_bulldozer_asm
PUBLIC cnv2_double_mainloop_sandybridge_asm PUBLIC cnv2_double_mainloop_sandybridge_asm
PUBLIC cnv2_rwz_mainloop_asm PUBLIC cnv2_rwz_mainloop_asm
PUBLIC cnv2_rwz_double_mainloop_asm PUBLIC cnv2_rwz_double_mainloop_asm
PUBLIC cnv2_upx_double_mainloop_zen3_asm
ALIGN(64) ALIGN(64)
cnv1_single_mainloop_asm PROC cnv1_single_mainloop_asm PROC

View File

@@ -8,7 +8,6 @@ PUBLIC cnv2_mainloop_bulldozer_asm
PUBLIC cnv2_double_mainloop_sandybridge_asm PUBLIC cnv2_double_mainloop_sandybridge_asm
PUBLIC cnv2_rwz_mainloop_asm PUBLIC cnv2_rwz_mainloop_asm
PUBLIC cnv2_rwz_double_mainloop_asm PUBLIC cnv2_rwz_double_mainloop_asm
PUBLIC cnv2_upx_double_mainloop_zen3_asm
ALIGN(64) ALIGN(64)
cnv1_single_mainloop_asm PROC cnv1_single_mainloop_asm PROC

File diff suppressed because it is too large Load Diff

View File

@@ -131,8 +131,8 @@ void JitCompilerA64::generateProgram(Program& program, ProgramConfiguration& con
// and w16, w10, ScratchpadL3Mask64 // and w16, w10, ScratchpadL3Mask64
emit32(0x121A0000 | 16 | (10 << 5) | ((RandomX_CurrentConfig.Log2_ScratchpadL3 - 7) << 10), code, codePos); emit32(0x121A0000 | 16 | (10 << 5) | ((RandomX_CurrentConfig.Log2_ScratchpadL3 - 7) << 10), code, codePos);
// and w17, w20, ScratchpadL3Mask64 // and w17, w18, ScratchpadL3Mask64
emit32(0x121A0000 | 17 | (20 << 5) | ((RandomX_CurrentConfig.Log2_ScratchpadL3 - 7) << 10), code, codePos); emit32(0x121A0000 | 17 | (18 << 5) | ((RandomX_CurrentConfig.Log2_ScratchpadL3 - 7) << 10), code, codePos);
codePos = PrologueSize; codePos = PrologueSize;
literalPos = ImulRcpLiteralsEnd; literalPos = ImulRcpLiteralsEnd;
@@ -148,16 +148,16 @@ void JitCompilerA64::generateProgram(Program& program, ProgramConfiguration& con
} }
// Update spMix2 // Update spMix2
// eor w20, config.readReg2, config.readReg3 // eor w18, config.readReg2, config.readReg3
emit32(ARMV8A::EOR32 | 20 | (IntRegMap[config.readReg2] << 5) | (IntRegMap[config.readReg3] << 16), code, codePos); emit32(ARMV8A::EOR32 | 18 | (IntRegMap[config.readReg2] << 5) | (IntRegMap[config.readReg3] << 16), code, codePos);
// Jump back to the main loop // Jump back to the main loop
const uint32_t offset = (((uint8_t*)randomx_program_aarch64_vm_instructions_end) - ((uint8_t*)randomx_program_aarch64)) - codePos; 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); 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)); 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 // and w10, w10, CacheLineAlignMask
codePos = (((uint8_t*)randomx_program_aarch64_cacheline_align_mask2) - ((uint8_t*)randomx_program_aarch64)); 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 // and w16, w10, ScratchpadL3Mask64
emit32(0x121A0000 | 16 | (10 << 5) | ((RandomX_CurrentConfig.Log2_ScratchpadL3 - 7) << 10), code, codePos); emit32(0x121A0000 | 16 | (10 << 5) | ((RandomX_CurrentConfig.Log2_ScratchpadL3 - 7) << 10), code, codePos);
// and w17, w20, ScratchpadL3Mask64 // and w17, w18, ScratchpadL3Mask64
emit32(0x121A0000 | 17 | (20 << 5) | ((RandomX_CurrentConfig.Log2_ScratchpadL3 - 7) << 10), code, codePos); emit32(0x121A0000 | 17 | (18 << 5) | ((RandomX_CurrentConfig.Log2_ScratchpadL3 - 7) << 10), code, codePos);
codePos = PrologueSize; codePos = PrologueSize;
literalPos = ImulRcpLiteralsEnd; literalPos = ImulRcpLiteralsEnd;
@@ -206,8 +206,8 @@ void JitCompilerA64::generateProgramLight(Program& program, ProgramConfiguration
} }
// Update spMix2 // Update spMix2
// eor w20, config.readReg2, config.readReg3 // eor w18, config.readReg2, config.readReg3
emit32(ARMV8A::EOR32 | 20 | (IntRegMap[config.readReg2] << 5) | (IntRegMap[config.readReg3] << 16), code, codePos); emit32(ARMV8A::EOR32 | 18 | (IntRegMap[config.readReg2] << 5) | (IntRegMap[config.readReg3] << 16), code, codePos);
// Jump back to the main loop // 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; 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 else
{ {
constexpr uint32_t tmp_reg = 20; constexpr uint32_t tmp_reg = 18;
emitMovImmediate(tmp_reg, imm, code, k); emitMovImmediate(tmp_reg, imm, code, k);
// add dst, src, tmp_reg // 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 k = codePos;
uint32_t imm = instr.getImm32(); 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); imm &= instr.getModMem() ? (RandomX_CurrentConfig.ScratchpadL1_Size - 1) : (RandomX_CurrentConfig.ScratchpadL2_Size - 1);
emitAddImmediate(tmp_reg, src, imm, code, k); 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 src = IntRegMap[instr.src];
const uint32_t dst = IntRegMap[instr.dst]; 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); emitMemLoad<tmp_reg>(dst, src, instr, code, k);
// add dst, dst, tmp_reg // 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 src = IntRegMap[instr.src];
const uint32_t dst = IntRegMap[instr.dst]; 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); emitMemLoad<tmp_reg>(dst, src, instr, code, k);
// sub dst, dst, tmp_reg // sub dst, dst, tmp_reg
@@ -637,7 +637,7 @@ void JitCompilerA64::h_IMUL_R(Instruction& instr, uint32_t& codePos)
if (src == dst) if (src == dst)
{ {
src = 20; src = 18;
emitMovImmediate(src, instr.getImm32(), code, k); 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 src = IntRegMap[instr.src];
const uint32_t dst = IntRegMap[instr.dst]; 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); emitMemLoad<tmp_reg>(dst, src, instr, code, k);
// sub dst, dst, tmp_reg // 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 src = IntRegMap[instr.src];
const uint32_t dst = IntRegMap[instr.dst]; 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); emitMemLoad<tmp_reg>(dst, src, instr, code, k);
// umulh dst, dst, tmp_reg // 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 src = IntRegMap[instr.src];
const uint32_t dst = IntRegMap[instr.dst]; 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); emitMemLoad<tmp_reg>(dst, src, instr, code, k);
// smulh dst, dst, tmp_reg // smulh dst, dst, tmp_reg
@@ -735,7 +735,7 @@ void JitCompilerA64::h_IMUL_RCP(Instruction& instr, uint32_t& codePos)
uint32_t k = codePos; uint32_t k = codePos;
constexpr uint32_t tmp_reg = 20; constexpr uint32_t tmp_reg = 18;
const uint32_t dst = IntRegMap[instr.dst]; const uint32_t dst = IntRegMap[instr.dst];
constexpr uint64_t N = 1ULL << 63; constexpr uint64_t N = 1ULL << 63;
@@ -754,9 +754,9 @@ void JitCompilerA64::h_IMUL_RCP(Instruction& instr, uint32_t& codePos)
literalPos -= sizeof(uint64_t); literalPos -= sizeof(uint64_t);
*(uint64_t*)(code + literalPos) = (q << shift) + ((r << shift) / divisor); *(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 // mul dst, dst, literal_reg
emit32(ARMV8A::MUL | dst | (dst << 5) | literal_regs[literal_id], code, k); 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) if (src == dst)
{ {
src = 20; src = 18;
emitMovImmediate(src, instr.getImm32(), code, k); 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 src = IntRegMap[instr.src];
const uint32_t dst = IntRegMap[instr.dst]; 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); emitMemLoad<tmp_reg>(dst, src, instr, code, k);
// eor dst, dst, tmp_reg // eor dst, dst, tmp_reg
@@ -850,7 +850,7 @@ void JitCompilerA64::h_IROL_R(Instruction& instr, uint32_t& codePos)
if (src != dst) if (src != dst)
{ {
constexpr uint32_t tmp_reg = 20; constexpr uint32_t tmp_reg = 18;
// sub tmp_reg, xzr, src // sub tmp_reg, xzr, src
emit32(ARMV8A::SUB | tmp_reg | (31 << 5) | (src << 16), code, k); 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; 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 | tmp_reg | (dst << 16), code, k);
emit32(ARMV8A::MOV_REG | dst | (src << 16), code, k); emit32(ARMV8A::MOV_REG | dst | (src << 16), code, k);
emit32(ARMV8A::MOV_REG | src | (tmp_reg << 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]; 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; constexpr uint32_t fpcr_tmp_reg = 8;
// ror tmp_reg, src, imm // 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 src = IntRegMap[instr.src];
const uint32_t dst = IntRegMap[instr.dst]; const uint32_t dst = IntRegMap[instr.dst];
constexpr uint32_t tmp_reg = 20; constexpr uint32_t tmp_reg = 18;
uint32_t imm = instr.getImm32(); uint32_t imm = instr.getImm32();

View File

@@ -72,9 +72,9 @@
# x15 -> "r7" # x15 -> "r7"
# x16 -> spAddr0 # x16 -> spAddr0
# x17 -> spAddr1 # x17 -> spAddr1
# x18 -> unused (platform register, don't touch it) # x18 -> temporary
# x19 -> temporary # x19 -> temporary
# x20 -> temporary # x20 -> literal for IMUL_RCP
# x21 -> literal for IMUL_RCP # x21 -> literal for IMUL_RCP
# x22 -> literal for IMUL_RCP # x22 -> literal for IMUL_RCP
# x23 -> literal for IMUL_RCP # x23 -> literal for IMUL_RCP
@@ -109,7 +109,7 @@ DECL(randomx_program_aarch64):
# Save callee-saved registers # Save callee-saved registers
sub sp, sp, 192 sub sp, sp, 192
stp x16, x17, [sp] stp x16, x17, [sp]
str x19, [sp, 16] stp x18, x19, [sp, 16]
stp x20, x21, [sp, 32] stp x20, x21, [sp, 32]
stp x22, x23, [sp, 48] stp x22, x23, [sp, 48]
stp x24, x25, [sp, 64] stp x24, x25, [sp, 64]
@@ -164,6 +164,7 @@ DECL(randomx_program_aarch64):
# Read literals # Read literals
ldr x0, literal_x0 ldr x0, literal_x0
ldr x11, literal_x11 ldr x11, literal_x11
ldr x20, literal_x20
ldr x21, literal_x21 ldr x21, literal_x21
ldr x22, literal_x22 ldr x22, literal_x22
ldr x23, literal_x23 ldr x23, literal_x23
@@ -195,11 +196,11 @@ DECL(randomx_program_aarch64):
DECL(randomx_program_aarch64_main_loop): DECL(randomx_program_aarch64_main_loop):
# spAddr0 = spMix1 & ScratchpadL3Mask64; # spAddr0 = spMix1 & ScratchpadL3Mask64;
# spAddr1 = (spMix1 >> 32) & ScratchpadL3Mask64; # spAddr1 = (spMix1 >> 32) & ScratchpadL3Mask64;
lsr x20, x10, 32 lsr x18, x10, 32
# Actual mask will be inserted by JIT compiler # Actual mask will be inserted by JIT compiler
and w16, w10, 1 and w16, w10, 1
and w17, w20, 1 and w17, w18, 1
# x16 = scratchpad + spAddr0 # x16 = scratchpad + spAddr0
# x17 = scratchpad + spAddr1 # x17 = scratchpad + spAddr1
@@ -207,31 +208,31 @@ DECL(randomx_program_aarch64_main_loop):
add x17, x17, x2 add x17, x17, x2
# xor integer registers with scratchpad data (spAddr0) # xor integer registers with scratchpad data (spAddr0)
ldp x20, x19, [x16] ldp x18, x19, [x16]
eor x4, x4, x20 eor x4, x4, x18
eor x5, x5, x19 eor x5, x5, x19
ldp x20, x19, [x16, 16] ldp x18, x19, [x16, 16]
eor x6, x6, x20 eor x6, x6, x18
eor x7, x7, x19 eor x7, x7, x19
ldp x20, x19, [x16, 32] ldp x18, x19, [x16, 32]
eor x12, x12, x20 eor x12, x12, x18
eor x13, x13, x19 eor x13, x13, x19
ldp x20, x19, [x16, 48] ldp x18, x19, [x16, 48]
eor x14, x14, x20 eor x14, x14, x18
eor x15, x15, x19 eor x15, x15, x19
# Load group F registers (spAddr1) # Load group F registers (spAddr1)
ldpsw x20, x19, [x17] ldpsw x18, x19, [x17]
ins v16.d[0], x20 ins v16.d[0], x18
ins v16.d[1], x19 ins v16.d[1], x19
ldpsw x20, x19, [x17, 8] ldpsw x18, x19, [x17, 8]
ins v17.d[0], x20 ins v17.d[0], x18
ins v17.d[1], x19 ins v17.d[1], x19
ldpsw x20, x19, [x17, 16] ldpsw x18, x19, [x17, 16]
ins v18.d[0], x20 ins v18.d[0], x18
ins v18.d[1], x19 ins v18.d[1], x19
ldpsw x20, x19, [x17, 24] ldpsw x18, x19, [x17, 24]
ins v19.d[0], x20 ins v19.d[0], x18
ins v19.d[1], x19 ins v19.d[1], x19
scvtf v16.2d, v16.2d scvtf v16.2d, v16.2d
scvtf v17.2d, v17.2d scvtf v17.2d, v17.2d
@@ -239,17 +240,17 @@ DECL(randomx_program_aarch64_main_loop):
scvtf v19.2d, v19.2d scvtf v19.2d, v19.2d
# Load group E registers (spAddr1) # Load group E registers (spAddr1)
ldpsw x20, x19, [x17, 32] ldpsw x18, x19, [x17, 32]
ins v20.d[0], x20 ins v20.d[0], x18
ins v20.d[1], x19 ins v20.d[1], x19
ldpsw x20, x19, [x17, 40] ldpsw x18, x19, [x17, 40]
ins v21.d[0], x20 ins v21.d[0], x18
ins v21.d[1], x19 ins v21.d[1], x19
ldpsw x20, x19, [x17, 48] ldpsw x18, x19, [x17, 48]
ins v22.d[0], x20 ins v22.d[0], x18
ins v22.d[1], x19 ins v22.d[1], x19
ldpsw x20, x19, [x17, 56] ldpsw x18, x19, [x17, 56]
ins v23.d[0], x20 ins v23.d[0], x18
ins v23.d[1], x19 ins v23.d[1], x19
scvtf v20.2d, v20.2d scvtf v20.2d, v20.2d
scvtf v21.2d, v21.2d scvtf v21.2d, v21.2d
@@ -272,6 +273,7 @@ DECL(randomx_program_aarch64_vm_instructions):
literal_x0: .fill 1,8,0 literal_x0: .fill 1,8,0
literal_x11: .fill 1,8,0 literal_x11: .fill 1,8,0
literal_x20: .fill 1,8,0
literal_x21: .fill 1,8,0 literal_x21: .fill 1,8,0
literal_x22: .fill 1,8,0 literal_x22: .fill 1,8,0
literal_x23: .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 lsr x10, x9, 32
# mx ^= r[readReg2] ^ r[readReg3]; # mx ^= r[readReg2] ^ r[readReg3];
eor x9, x9, x20 eor x9, x9, x18
# Calculate dataset pointer for dataset prefetch # Calculate dataset pointer for dataset prefetch
mov w20, w9 mov w18, w9
DECL(randomx_program_aarch64_cacheline_align_mask1): DECL(randomx_program_aarch64_cacheline_align_mask1):
# Actual mask will be inserted by JIT compiler # Actual mask will be inserted by JIT compiler
and x20, x20, 1 and x18, x18, 1
add x20, x20, x1 add x18, x18, x1
# Prefetch dataset data # Prefetch dataset data
prfm pldl2strm, [x20] prfm pldl2strm, [x18]
# mx <-> ma # mx <-> ma
ror x9, x9, 32 ror x9, x9, 32
@@ -329,17 +331,17 @@ DECL(randomx_program_aarch64_cacheline_align_mask2):
DECL(randomx_program_aarch64_xor_with_dataset_line): DECL(randomx_program_aarch64_xor_with_dataset_line):
# xor integer registers with dataset data # xor integer registers with dataset data
ldp x20, x19, [x10] ldp x18, x19, [x10]
eor x4, x4, x20 eor x4, x4, x18
eor x5, x5, x19 eor x5, x5, x19
ldp x20, x19, [x10, 16] ldp x18, x19, [x10, 16]
eor x6, x6, x20 eor x6, x6, x18
eor x7, x7, x19 eor x7, x7, x19
ldp x20, x19, [x10, 32] ldp x18, x19, [x10, 32]
eor x12, x12, x20 eor x12, x12, x18
eor x13, x13, x19 eor x13, x13, x19
ldp x20, x19, [x10, 48] ldp x18, x19, [x10, 48]
eor x14, x14, x20 eor x14, x14, x18
eor x15, x15, x19 eor x15, x15, x19
DECL(randomx_program_aarch64_update_spMix1): DECL(randomx_program_aarch64_update_spMix1):
@@ -382,7 +384,7 @@ DECL(randomx_program_aarch64_update_spMix1):
# Restore callee-saved registers # Restore callee-saved registers
ldp x16, x17, [sp] ldp x16, x17, [sp]
ldr x19, [sp, 16] ldp x18, x19, [sp, 16]
ldp x20, x21, [sp, 32] ldp x20, x21, [sp, 32]
ldp x22, x23, [sp, 48] ldp x22, x23, [sp, 48]
ldp x24, x25, [sp, 64] ldp x24, x25, [sp, 64]
@@ -403,7 +405,7 @@ DECL(randomx_program_aarch64_vm_instructions_end_light):
stp x2, x30, [sp, 80] stp x2, x30, [sp, 80]
# mx ^= r[readReg2] ^ r[readReg3]; # mx ^= r[readReg2] ^ r[readReg3];
eor x9, x9, x20 eor x9, x9, x18
# mx <-> ma # mx <-> ma
ror x9, x9, 32 ror x9, x9, 32
@@ -445,8 +447,8 @@ DECL(randomx_program_aarch64_light_dataset_offset):
# x3 -> end item # x3 -> end item
DECL(randomx_init_dataset_aarch64): DECL(randomx_init_dataset_aarch64):
# Save x20 (used as temporary, but must be saved to not break ABI) and x30 (return address) # Save x30 (return address)
stp x20, x30, [sp, -16]! str x30, [sp, -16]!
# Load pointer to cache memory # Load pointer to cache memory
ldr x0, [x0] ldr x0, [x0]
@@ -458,8 +460,8 @@ DECL(randomx_init_dataset_aarch64_main_loop):
cmp x2, x3 cmp x2, x3
bne DECL(randomx_init_dataset_aarch64_main_loop) bne DECL(randomx_init_dataset_aarch64_main_loop)
# Restore x20 and x30 # Restore x30 (return address)
ldp x20, x30, [sp], 16 ldr x30, [sp], 16
ret ret

View File

@@ -1,6 +1,6 @@
/* XMRig /* XMRig
* Copyright (c) 2018-2024 SChernykh <https://github.com/SChernykh> * Copyright (c) 2018-2023 SChernykh <https://github.com/SChernykh>
* Copyright (c) 2016-2024 XMRig <https://github.com/xmrig>, <support@xmrig.com> * Copyright (c) 2016-2023 XMRig <https://github.com/xmrig>, <support@xmrig.com>
* *
* This program is free software: you can redistribute it and/or modify * 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 * it under the terms of the GNU General Public License as published by
@@ -22,14 +22,14 @@
#define APP_ID "xmrig" #define APP_ID "xmrig"
#define APP_NAME "XMRig" #define APP_NAME "XMRig"
#define APP_DESC "XMRig miner" #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_DOMAIN "xmrig.com"
#define APP_SITE "www.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_KIND "miner"
#define APP_VER_MAJOR 6 #define APP_VER_MAJOR 6
#define APP_VER_MINOR 21 #define APP_VER_MINOR 20
#define APP_VER_PATCH 1 #define APP_VER_PATCH 1
#ifdef _MSC_VER #ifdef _MSC_VER