mirror of
https://github.com/xmrig/xmrig.git
synced 2025-12-10 01:02:40 -05:00
Compare commits
23 Commits
8aa110d182
...
v6.21.1
| Author | SHA1 | Date | |
|---|---|---|---|
|
|
a5aa2c9042 | ||
|
|
fa35a32eee | ||
|
|
7b6ce59821 | ||
|
|
33315ba2ef | ||
|
|
2c9c40d623 | ||
|
|
daa6328418 | ||
|
|
8afd4d5f2f | ||
|
|
77e2f3a028 | ||
|
|
206295c6cb | ||
|
|
07e1e77c4f | ||
|
|
50a98a4bb1 | ||
|
|
c50369d65d | ||
|
|
592b0c9c76 | ||
|
|
89eab0eff2 | ||
|
|
8084ff37a5 | ||
|
|
7cf3db7750 | ||
|
|
4bda6e054d | ||
|
|
64a0ed413b | ||
|
|
0b59b7eb43 | ||
|
|
ae6b10b5a4 | ||
|
|
705a7eac0c | ||
|
|
10bfffe033 | ||
|
|
4131aa4754 |
15
CHANGELOG.md
15
CHANGELOG.md
@@ -1,3 +1,18 @@
|
||||
# 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`.
|
||||
|
||||
@@ -138,6 +138,93 @@ __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
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@@ -0,0 +1,58 @@
|
||||
/* 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, >hreads, <hreads);
|
||||
}
|
||||
|
||||
|
||||
// __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);
|
||||
}
|
||||
@@ -0,0 +1,50 @@
|
||||
/* XMRig
|
||||
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||
* Copyright 2018-2019 SChernykh <https://github.com/SChernykh>
|
||||
* Copyright 2016-2019 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||
*
|
||||
* This program is free software: you can redistribute it and/or modify
|
||||
* it under the terms of the GNU General Public License as published by
|
||||
* the Free Software Foundation, either version 3 of the License, or
|
||||
* (at your option) any later version.
|
||||
*
|
||||
* This program is distributed in the hope that it will be useful,
|
||||
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
* GNU General Public License for more details.
|
||||
*
|
||||
* You should have received a copy of the GNU General Public License
|
||||
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||
*/
|
||||
|
||||
#ifndef XMRIG_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 */
|
||||
@@ -80,6 +80,7 @@ if (WITH_OPENCL)
|
||||
if (WITH_RANDOMX)
|
||||
list(APPEND HEADERS_BACKEND_OPENCL
|
||||
src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h
|
||||
src/backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.h
|
||||
src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.h
|
||||
src/backend/opencl/kernels/rx/ExecuteVmKernel.h
|
||||
src/backend/opencl/kernels/rx/FillAesKernel.h
|
||||
@@ -96,6 +97,7 @@ if (WITH_OPENCL)
|
||||
list(APPEND SOURCES_BACKEND_OPENCL
|
||||
src/backend/opencl/generators/ocl_generic_rx_generator.cpp
|
||||
src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.cpp
|
||||
src/backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.cpp
|
||||
src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.cpp
|
||||
src/backend/opencl/kernels/rx/ExecuteVmKernel.cpp
|
||||
src/backend/opencl/kernels/rx/FillAesKernel.cpp
|
||||
|
||||
@@ -25,6 +25,7 @@
|
||||
#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"
|
||||
@@ -71,6 +72,7 @@ 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;
|
||||
@@ -87,12 +89,28 @@ void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
|
||||
{
|
||||
static const uint32_t zero = 0;
|
||||
|
||||
m_blake2b_initial_hash->setNonce(nonce);
|
||||
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_find_shares->setNonce(nonce);
|
||||
|
||||
enqueueWriteBuffer(m_output, CL_FALSE, sizeof(cl_uint) * 0xFF, sizeof(uint32_t), &zero);
|
||||
|
||||
m_blake2b_initial_hash->enqueue(m_queue, m_intensity);
|
||||
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_fillAes1Rx4_scratchpad->enqueue(m_queue, m_intensity);
|
||||
|
||||
const uint32_t programCount = RxAlgo::programCount(m_algorithm);
|
||||
@@ -134,7 +152,11 @@ 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());
|
||||
}
|
||||
|
||||
@@ -166,6 +188,9 @@ 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");
|
||||
|
||||
|
||||
@@ -35,6 +35,7 @@ namespace xmrig {
|
||||
|
||||
class Blake2bHashRegistersKernel;
|
||||
class Blake2bInitialHashKernel;
|
||||
class Blake2bInitialHashDoubleKernel;
|
||||
class FillAesKernel;
|
||||
class FindSharesKernel;
|
||||
class HashAesKernel;
|
||||
@@ -58,21 +59,24 @@ 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;
|
||||
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;
|
||||
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;
|
||||
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;
|
||||
};
|
||||
|
||||
|
||||
|
||||
@@ -39,6 +39,7 @@
|
||||
|
||||
|
||||
#include <thread>
|
||||
#include <iostream>
|
||||
|
||||
|
||||
namespace xmrig {
|
||||
@@ -80,7 +81,8 @@ static rapidjson::Value getResources(rapidjson::Document &doc)
|
||||
|
||||
xmrig::Api::Api(Base *base) :
|
||||
m_base(base),
|
||||
m_timestamp(Chrono::currentMSecsSinceEpoch())
|
||||
m_timestamp(Chrono::currentMSecsSinceEpoch()),
|
||||
m_httpd(nullptr)
|
||||
{
|
||||
base->addListener(this);
|
||||
|
||||
@@ -91,7 +93,11 @@ xmrig::Api::Api(Base *base) :
|
||||
xmrig::Api::~Api()
|
||||
{
|
||||
# ifdef XMRIG_FEATURE_HTTP
|
||||
delete m_httpd;
|
||||
if (m_httpd) {
|
||||
m_httpd->stop();
|
||||
delete m_httpd;
|
||||
m_httpd = nullptr; // Ensure the pointer is set to nullptr after deletion
|
||||
}
|
||||
# endif
|
||||
}
|
||||
|
||||
@@ -109,8 +115,14 @@ void xmrig::Api::start()
|
||||
genWorkerId(m_base->config()->apiWorkerId());
|
||||
|
||||
# ifdef XMRIG_FEATURE_HTTP
|
||||
m_httpd = new Httpd(m_base);
|
||||
m_httpd->start();
|
||||
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;
|
||||
}
|
||||
}
|
||||
# endif
|
||||
}
|
||||
|
||||
@@ -118,7 +130,9 @@ void xmrig::Api::start()
|
||||
void xmrig::Api::stop()
|
||||
{
|
||||
# ifdef XMRIG_FEATURE_HTTP
|
||||
m_httpd->stop();
|
||||
if (m_httpd) {
|
||||
m_httpd->stop();
|
||||
}
|
||||
# endif
|
||||
}
|
||||
|
||||
@@ -126,13 +140,15 @@ void xmrig::Api::stop()
|
||||
void xmrig::Api::tick()
|
||||
{
|
||||
# ifdef XMRIG_FEATURE_HTTP
|
||||
if (m_httpd->isBound() || !m_base->config()->http().isEnabled()) {
|
||||
if (!m_httpd || !m_base->config()->http().isEnabled() || m_httpd->isBound()) {
|
||||
return;
|
||||
}
|
||||
|
||||
if (++m_ticks % 10 == 0) {
|
||||
m_ticks = 0;
|
||||
m_httpd->start();
|
||||
if (m_httpd) {
|
||||
m_httpd->start();
|
||||
}
|
||||
}
|
||||
# endif
|
||||
}
|
||||
|
||||
@@ -54,6 +54,7 @@ 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 ") },
|
||||
};
|
||||
|
||||
|
||||
|
||||
@@ -40,6 +40,7 @@ public:
|
||||
RAVEN,
|
||||
WOWNERO,
|
||||
ZEPHYR,
|
||||
TOWNFORGE,
|
||||
MAX
|
||||
};
|
||||
|
||||
|
||||
@@ -1,7 +1,7 @@
|
||||
/* XMRig
|
||||
* Copyright (c) 2019 Howard Chu <https://github.com/hyc>
|
||||
* Copyright (c) 2018-2021 SChernykh <https://github.com/SChernykh>
|
||||
* Copyright (c) 2016-2021 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||
* Copyright (c) 2018-2024 SChernykh <https://github.com/SChernykh>
|
||||
* Copyright (c) 2016-2024 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||
*
|
||||
* This program is free software: you can redistribute it and/or modify
|
||||
* it under the terms of the GNU General Public License as published by
|
||||
@@ -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,11 @@ public:
|
||||
static const char *kSOCKS5;
|
||||
static const char *kSubmitToOrigin;
|
||||
static const char *kTls;
|
||||
static const char* kSni;
|
||||
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;
|
||||
|
||||
@@ -1,8 +1,8 @@
|
||||
/* XMRig
|
||||
* Copyright (c) 2012-2013 The Cryptonote developers
|
||||
* Copyright (c) 2014-2021 The Monero Project
|
||||
* Copyright (c) 2018-2021 SChernykh <https://github.com/SChernykh>
|
||||
* Copyright (c) 2016-2021 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
|
||||
@@ -198,7 +198,7 @@ bool xmrig::BlockTemplate::parse(bool hashes)
|
||||
}
|
||||
|
||||
if (m_coin == Coin::ZEPHYR) {
|
||||
uint8_t pricing_record[24];
|
||||
uint8_t pricing_record[120];
|
||||
ar(pricing_record);
|
||||
}
|
||||
|
||||
@@ -207,7 +207,11 @@ bool xmrig::BlockTemplate::parse(bool hashes)
|
||||
setOffset(MINER_TX_PREFIX_OFFSET, ar.index());
|
||||
|
||||
ar(m_txVersion);
|
||||
ar(m_unlockTime);
|
||||
|
||||
if (m_coin != Coin::TOWNFORGE) {
|
||||
ar(m_unlockTime);
|
||||
}
|
||||
|
||||
ar(m_numInputs);
|
||||
|
||||
// must be 1 input
|
||||
@@ -225,8 +229,12 @@ bool xmrig::BlockTemplate::parse(bool hashes)
|
||||
ar(m_height);
|
||||
ar(m_numOutputs);
|
||||
|
||||
const uint64_t expected_outputs = (m_coin == Coin::ZEPHYR) ? 2 : 1;
|
||||
if (m_numOutputs != expected_outputs) {
|
||||
if (m_coin == Coin::ZEPHYR) {
|
||||
if (m_numOutputs < 2) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
else if (m_numOutputs != 1) {
|
||||
return false;
|
||||
}
|
||||
|
||||
@@ -252,28 +260,34 @@ bool xmrig::BlockTemplate::parse(bool hashes)
|
||||
ar.skip(asset_type_len);
|
||||
ar(m_viewTag);
|
||||
|
||||
uint64_t amount2;
|
||||
ar(amount2);
|
||||
for (uint64_t k = 1; k < m_numOutputs; ++k) {
|
||||
uint64_t amount2;
|
||||
ar(amount2);
|
||||
|
||||
uint8_t output_type2;
|
||||
ar(output_type2);
|
||||
if (output_type2 != 2) {
|
||||
return false;
|
||||
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);
|
||||
}
|
||||
|
||||
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());
|
||||
@@ -329,6 +343,11 @@ 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;
|
||||
|
||||
@@ -1,8 +1,8 @@
|
||||
/* XMRig
|
||||
* Copyright (c) 2012-2013 The Cryptonote developers
|
||||
* Copyright (c) 2014-2021 The Monero Project
|
||||
* Copyright (c) 2018-2021 SChernykh <https://github.com/SChernykh>
|
||||
* Copyright (c) 2016-2021 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
|
||||
@@ -33,6 +33,25 @@
|
||||
|
||||
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;
|
||||
@@ -114,6 +133,10 @@ 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;
|
||||
}
|
||||
}
|
||||
@@ -228,6 +251,16 @@ 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);
|
||||
|
||||
@@ -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-2021 SChernykh <https://github.com/SChernykh>
|
||||
* Copyright (c) 2016-2021 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||
* Copyright (c) 2018-2024 SChernykh <https://github.com/SChernykh>
|
||||
* Copyright (c) 2016-2024 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||
*
|
||||
* This program is free software: you can redistribute it and/or modify
|
||||
* it under the terms of the GNU General Public License as published by
|
||||
@@ -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 daemon's zmq-pub port number (only use it if daemon has it enabled)\n";
|
||||
u += " --daemon-zmq-port=N 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";
|
||||
|
||||
@@ -37,36 +37,14 @@ class CnAlgo
|
||||
public:
|
||||
constexpr CnAlgo() {};
|
||||
|
||||
# define ASSERT_CN static_assert(Algorithm::isCN(ALGO), "invalid CRYPTONIGHT algorithm")
|
||||
constexpr inline Algorithm::Id base() const { ASSERT_CN; return Algorithm::base(ALGO); }
|
||||
constexpr inline size_t memory() const { ASSERT_CN; return Algorithm::l3(ALGO); }
|
||||
constexpr inline uint32_t iterations() const { ASSERT_CN; return CN_ITER; }
|
||||
constexpr inline Algorithm::Id base() const { static_assert(Algorithm::isCN(ALGO), "invalid CRYPTONIGHT algorithm"); return Algorithm::base(ALGO); }
|
||||
constexpr inline bool isHeavy() const { return Algorithm::family(ALGO) == Algorithm::CN_HEAVY; }
|
||||
constexpr inline bool isR() const { return ALGO == Algorithm::CN_R; }
|
||||
constexpr inline size_t memory() const { static_assert(Algorithm::isCN(ALGO), "invalid CRYPTONIGHT algorithm"); return Algorithm::l3(ALGO); }
|
||||
constexpr inline uint32_t iterations() const { static_assert(Algorithm::isCN(ALGO), "invalid CRYPTONIGHT algorithm"); return CN_ITER; }
|
||||
constexpr inline uint32_t mask() const { return static_cast<uint32_t>(((memory() - 1) / 16) * 16); }
|
||||
constexpr inline uint32_t half_mem() const { return mask() < memory() / 2; }
|
||||
|
||||
constexpr inline bool isBase1() const { ASSERT_CN; return Algorithm::base(ALGO) == Algorithm::CN_1; }
|
||||
constexpr inline bool isBase2() const { ASSERT_CN; return Algorithm::base(ALGO) == Algorithm::CN_2; }
|
||||
constexpr inline bool is2() const { return ALGO == Algorithm::CN_2; }
|
||||
constexpr inline bool isR() const { return ALGO == Algorithm::CN_R; }
|
||||
constexpr inline bool isHalf() const { return ALGO == Algorithm::CN_HALF; }
|
||||
constexpr inline bool isRTO() const { return ALGO == Algorithm::CN_RTO; }
|
||||
constexpr inline bool isRWZ() const { return ALGO == Algorithm::CN_RWZ; }
|
||||
constexpr inline bool isZLS() const { return ALGO == Algorithm::CN_ZLS; }
|
||||
constexpr inline bool isDouble() const { return ALGO == Algorithm::CN_DOUBLE; }
|
||||
constexpr inline bool isCCX() const { return ALGO == Algorithm::CN_CCX; }
|
||||
constexpr inline bool isHeavy() const { ASSERT_CN; return Algorithm::family(ALGO) == Algorithm::CN_HEAVY; }
|
||||
constexpr inline bool isHeavyTube() const { return ALGO == Algorithm::CN_HEAVY_TUBE; }
|
||||
constexpr inline bool isHeavyXHV() const { return ALGO == Algorithm::CN_HEAVY_XHV; }
|
||||
constexpr inline bool isPico0() const { return ALGO == Algorithm::CN_PICO_0; }
|
||||
constexpr inline bool isPicoTLO() const { return ALGO == Algorithm::CN_PICO_TLO; }
|
||||
constexpr inline bool isUPX2() const { return ALGO == Algorithm::CN_UPX2; }
|
||||
constexpr inline bool isGR0() const { return ALGO == Algorithm::CN_GR_0; }
|
||||
constexpr inline bool isGR1() const { return ALGO == Algorithm::CN_GR_1; }
|
||||
constexpr inline bool isGR2() const { return ALGO == Algorithm::CN_GR_2; }
|
||||
constexpr inline bool isGR3() const { return ALGO == Algorithm::CN_GR_3; }
|
||||
constexpr inline bool isGR4() const { return ALGO == Algorithm::CN_GR_4; }
|
||||
constexpr inline bool isGR5() const { return ALGO == Algorithm::CN_GR_5; }
|
||||
|
||||
inline static uint32_t iterations(Algorithm::Id algo)
|
||||
{
|
||||
switch (algo) {
|
||||
|
||||
@@ -603,7 +603,7 @@ static inline void cryptonight_monero_tweak(uint64_t *mem_out, const uint8_t *l,
|
||||
constexpr CnAlgo<ALGO> props;
|
||||
|
||||
if (props.base() == Algorithm::CN_2) {
|
||||
VARIANT2_SHUFFLE(l, idx, ax0, bx0, bx1, cx, ((props.isRWZ() || props.isUPX2()) ? 1 : 0));
|
||||
VARIANT2_SHUFFLE(l, idx, ax0, bx0, bx1, cx, (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? 1 : 0));
|
||||
_mm_store_si128(reinterpret_cast<__m128i *>(mem_out), _mm_xor_si128(bx0, cx));
|
||||
} else {
|
||||
__m128i tmp = _mm_xor_si128(bx0, cx);
|
||||
@@ -665,8 +665,15 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
|
||||
|
||||
constexpr CnAlgo<ALGO> props;
|
||||
constexpr size_t MASK = props.mask();
|
||||
constexpr Algorithm::Id BASE = props.base();
|
||||
|
||||
if (props.isBase1() && size < 43) {
|
||||
# ifdef XMRIG_ALGO_CN_HEAVY
|
||||
constexpr bool IS_CN_HEAVY_TUBE = ALGO == Algorithm::CN_HEAVY_TUBE;
|
||||
# else
|
||||
constexpr bool IS_CN_HEAVY_TUBE = false;
|
||||
# endif
|
||||
|
||||
if (BASE == Algorithm::CN_1 && size < 43) {
|
||||
memset(output, 0, 32);
|
||||
return;
|
||||
}
|
||||
@@ -687,7 +694,10 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
|
||||
V4_Instruction code[256];
|
||||
const int code_size = v4_random_math_init<ALGO>(code, height);
|
||||
|
||||
v4_soft_aes_compile_code(code, code_size, reinterpret_cast<void*>(ctx[0]->generated_code), Assembly::NONE);
|
||||
if (ALGO == Algorithm::CN_R) {
|
||||
v4_soft_aes_compile_code(code, code_size, reinterpret_cast<void*>(ctx[0]->generated_code), Assembly::NONE);
|
||||
}
|
||||
|
||||
ctx[0]->generated_code_data = { ALGO, height };
|
||||
}
|
||||
|
||||
@@ -708,26 +718,26 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
|
||||
__m128i bx1 = _mm_set_epi64x(static_cast<int64_t>(h0[9] ^ h0[11]), static_cast<int64_t>(h0[8] ^ h0[10]));
|
||||
|
||||
__m128 conc_var;
|
||||
if (props.isCCX()) {
|
||||
if (ALGO == Algorithm::CN_CCX) {
|
||||
conc_var = _mm_setzero_ps();
|
||||
RESTORE_ROUNDING_MODE();
|
||||
}
|
||||
|
||||
for (size_t i = 0; i < props.iterations(); i++) {
|
||||
__m128i cx;
|
||||
if (props.isHeavyTube() || !SOFT_AES) {
|
||||
if (IS_CN_HEAVY_TUBE || !SOFT_AES) {
|
||||
cx = _mm_load_si128(reinterpret_cast<const __m128i *>(&l0[interleaved_index<interleave>(idx0 & MASK)]));
|
||||
if (props.isCCX()) {
|
||||
if (ALGO == Algorithm::CN_CCX) {
|
||||
cryptonight_conceal_tweak(cx, conc_var);
|
||||
}
|
||||
}
|
||||
|
||||
const __m128i ax0 = _mm_set_epi64x(static_cast<int64_t>(ah0), static_cast<int64_t>(al0));
|
||||
if (props.isHeavyTube()) {
|
||||
if (IS_CN_HEAVY_TUBE) {
|
||||
cx = aes_round_tweak_div(cx, ax0);
|
||||
}
|
||||
else if (SOFT_AES) {
|
||||
if (props.isCCX()) {
|
||||
if (ALGO == Algorithm::CN_CCX) {
|
||||
cx = _mm_load_si128(reinterpret_cast<const __m128i*>(&l0[interleaved_index<interleave>(idx0 & MASK)]));
|
||||
cryptonight_conceal_tweak(cx, conc_var);
|
||||
cx = soft_aesenc(&cx, ax0, reinterpret_cast<const uint32_t*>(saes_table));
|
||||
@@ -740,7 +750,7 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
|
||||
cx = _mm_aesenc_si128(cx, ax0);
|
||||
}
|
||||
|
||||
if (props.isBase1() || props.isBase2()) {
|
||||
if (BASE == Algorithm::CN_1 || BASE == Algorithm::CN_2) {
|
||||
cryptonight_monero_tweak<ALGO>(reinterpret_cast<uint64_t*>(&l0[interleaved_index<interleave>(idx0 & MASK)]), l0, idx0 & MASK, ax0, bx0, bx1, cx);
|
||||
} else {
|
||||
_mm_store_si128(reinterpret_cast<__m128i *>(&l0[interleaved_index<interleave>(idx0 & MASK)]), _mm_xor_si128(bx0, cx));
|
||||
@@ -752,11 +762,13 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
|
||||
cl = (reinterpret_cast<uint64_t*>(&l0[interleaved_index<interleave>(idx0 & MASK)]))[0];
|
||||
ch = (reinterpret_cast<uint64_t*>(&l0[interleaved_index<interleave>(idx0 & MASK)]))[1];
|
||||
|
||||
if (props.isBase2()) {
|
||||
if (BASE == Algorithm::CN_2) {
|
||||
if (props.isR()) {
|
||||
VARIANT4_RANDOM_MATH(0, al0, ah0, cl, bx0, bx1);
|
||||
al0 ^= r0[2] | (static_cast<uint64_t>(r0[3]) << 32);
|
||||
ah0 ^= r0[0] | (static_cast<uint64_t>(r0[1]) << 32);
|
||||
if (ALGO == Algorithm::CN_R) {
|
||||
al0 ^= r0[2] | (static_cast<uint64_t>(r0[3]) << 32);
|
||||
ah0 ^= r0[0] | (static_cast<uint64_t>(r0[1]) << 32);
|
||||
}
|
||||
} else {
|
||||
VARIANT2_INTEGER_MATH(0, cl, cx);
|
||||
}
|
||||
@@ -764,11 +776,11 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
|
||||
|
||||
lo = __umul128(idx0, cl, &hi);
|
||||
|
||||
if (props.isBase2()) {
|
||||
if (props.isR()) {
|
||||
if (BASE == Algorithm::CN_2) {
|
||||
if (ALGO == Algorithm::CN_R) {
|
||||
VARIANT2_SHUFFLE(l0, idx0 & MASK, ax0, bx0, bx1, cx, 0);
|
||||
} else {
|
||||
VARIANT2_SHUFFLE2(l0, idx0 & MASK, ax0, bx0, bx1, hi, lo, ((props.isRWZ() || props.isUPX2()) ? 1 : 0));
|
||||
VARIANT2_SHUFFLE2(l0, idx0 & MASK, ax0, bx0, bx1, hi, lo, (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? 1 : 0));
|
||||
}
|
||||
}
|
||||
|
||||
@@ -777,9 +789,9 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
|
||||
|
||||
reinterpret_cast<uint64_t*>(&l0[interleaved_index<interleave>(idx0 & MASK)])[0] = al0;
|
||||
|
||||
if (props.isHeavyTube() || props.isRTO()) {
|
||||
if (IS_CN_HEAVY_TUBE || ALGO == Algorithm::CN_RTO) {
|
||||
reinterpret_cast<uint64_t*>(&l0[interleaved_index<interleave>(idx0 & MASK)])[1] = ah0 ^ tweak1_2_0 ^ al0;
|
||||
} else if (props.isBase1()) {
|
||||
} else if (BASE == Algorithm::CN_1) {
|
||||
reinterpret_cast<uint64_t*>(&l0[interleaved_index<interleave>(idx0 & MASK)])[1] = ah0 ^ tweak1_2_0;
|
||||
} else {
|
||||
reinterpret_cast<uint64_t*>(&l0[interleaved_index<interleave>(idx0 & MASK)])[1] = ah0;
|
||||
@@ -807,7 +819,7 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
|
||||
|
||||
((int64_t*)&l0[interleaved_index<interleave>(idx0 & MASK)])[0] = n ^ q;
|
||||
|
||||
if (props.isHeavyXHV()) {
|
||||
if (ALGO == Algorithm::CN_HEAVY_XHV) {
|
||||
d = ~d;
|
||||
}
|
||||
|
||||
@@ -815,7 +827,7 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
|
||||
}
|
||||
# endif
|
||||
|
||||
if (props.isBase2()) {
|
||||
if (BASE == Algorithm::CN_2) {
|
||||
bx1 = bx0;
|
||||
}
|
||||
|
||||
@@ -948,7 +960,7 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
|
||||
}
|
||||
cn_explode_scratchpad<ALGO, false, 0>(ctx[0]);
|
||||
|
||||
if (props.is2()) {
|
||||
if (ALGO == Algorithm::CN_2) {
|
||||
if (ASM == Assembly::INTEL) {
|
||||
cnv2_mainloop_ivybridge_asm(ctx);
|
||||
}
|
||||
@@ -959,7 +971,7 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
|
||||
cnv2_mainloop_bulldozer_asm(ctx);
|
||||
}
|
||||
}
|
||||
else if (props.isHalf()) {
|
||||
else if (ALGO == Algorithm::CN_HALF) {
|
||||
if (ASM == Assembly::INTEL) {
|
||||
cn_half_mainloop_ivybridge_asm(ctx);
|
||||
}
|
||||
@@ -971,7 +983,7 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
|
||||
}
|
||||
}
|
||||
# ifdef XMRIG_ALGO_CN_PICO
|
||||
else if (props.isPico0()) {
|
||||
else if (ALGO == Algorithm::CN_PICO_0) {
|
||||
if (ASM == Assembly::INTEL) {
|
||||
cn_trtl_mainloop_ivybridge_asm(ctx);
|
||||
}
|
||||
@@ -982,7 +994,7 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
|
||||
cn_trtl_mainloop_bulldozer_asm(ctx);
|
||||
}
|
||||
}
|
||||
else if (props.isPicoTLO()) {
|
||||
else if (ALGO == Algorithm::CN_PICO_TLO) {
|
||||
if (ASM == Assembly::INTEL) {
|
||||
cn_tlo_mainloop_ivybridge_asm(ctx);
|
||||
}
|
||||
@@ -994,10 +1006,10 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
|
||||
}
|
||||
}
|
||||
# endif
|
||||
else if (props.isRWZ()) {
|
||||
else if (ALGO == Algorithm::CN_RWZ) {
|
||||
cnv2_rwz_mainloop_asm(ctx);
|
||||
}
|
||||
else if (props.isZLS()) {
|
||||
else if (ALGO == Algorithm::CN_ZLS) {
|
||||
if (ASM == Assembly::INTEL) {
|
||||
cn_zls_mainloop_ivybridge_asm(ctx);
|
||||
}
|
||||
@@ -1008,7 +1020,7 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
|
||||
cn_zls_mainloop_bulldozer_asm(ctx);
|
||||
}
|
||||
}
|
||||
else if (props.isDouble()) {
|
||||
else if (ALGO == Algorithm::CN_DOUBLE) {
|
||||
if (ASM == Assembly::INTEL) {
|
||||
cn_double_mainloop_ivybridge_asm(ctx);
|
||||
}
|
||||
@@ -1020,7 +1032,7 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
|
||||
}
|
||||
}
|
||||
# ifdef XMRIG_ALGO_CN_FEMTO
|
||||
else if (props.isUPX2()) {
|
||||
else if (ALGO == Algorithm::CN_UPX2) {
|
||||
cn_upx2_mainloop_asm(ctx);
|
||||
}
|
||||
# endif
|
||||
@@ -1066,22 +1078,22 @@ inline void cryptonight_double_hash_asm(const uint8_t *__restrict__ input, size_
|
||||
cn_explode_scratchpad<ALGO, false, 0>(ctx[1]);
|
||||
}
|
||||
|
||||
if (props.is2()) {
|
||||
if (ALGO == Algorithm::CN_2) {
|
||||
cnv2_double_mainloop_sandybridge_asm(ctx);
|
||||
}
|
||||
else if (props.isHalf()){
|
||||
else if (ALGO == Algorithm::CN_HALF) {
|
||||
cn_half_double_mainloop_sandybridge_asm(ctx);
|
||||
}
|
||||
# ifdef XMRIG_ALGO_CN_PICO
|
||||
else if (props.isPico0()) {
|
||||
else if (ALGO == Algorithm::CN_PICO_0) {
|
||||
cn_trtl_double_mainloop_sandybridge_asm(ctx);
|
||||
}
|
||||
else if (props.isPicoTLO()) {
|
||||
else if (ALGO == Algorithm::CN_PICO_TLO) {
|
||||
cn_tlo_double_mainloop_sandybridge_asm(ctx);
|
||||
}
|
||||
# endif
|
||||
# ifdef XMRIG_ALGO_CN_FEMTO
|
||||
else if (props.isUPX2()) {
|
||||
else if (ALGO == Algorithm::CN_UPX2) {
|
||||
if (Cpu::info()->arch() == ICpuInfo::ARCH_ZEN3) {
|
||||
cnv2_upx_double_mainloop_zen3_asm(ctx);
|
||||
}
|
||||
@@ -1090,13 +1102,13 @@ inline void cryptonight_double_hash_asm(const uint8_t *__restrict__ input, size_
|
||||
}
|
||||
}
|
||||
# endif
|
||||
else if (props.isRWZ()) {
|
||||
else if (ALGO == Algorithm::CN_RWZ) {
|
||||
cnv2_rwz_double_mainloop_asm(ctx);
|
||||
}
|
||||
else if (props.isZLS()) {
|
||||
else if (ALGO == Algorithm::CN_ZLS) {
|
||||
cn_zls_double_mainloop_sandybridge_asm(ctx);
|
||||
}
|
||||
else if (props.isDouble()) {
|
||||
else if (ALGO == Algorithm::CN_DOUBLE) {
|
||||
cn_double_double_mainloop_sandybridge_asm(ctx);
|
||||
}
|
||||
else if (props.isR()) {
|
||||
@@ -1134,8 +1146,9 @@ template<Algorithm::Id ALGO>
|
||||
static NOINLINE void cryptonight_single_hash_gr_sse41(const uint8_t* __restrict__ input, size_t size, uint8_t* __restrict__ output, cryptonight_ctx** __restrict__ ctx, uint64_t height)
|
||||
{
|
||||
constexpr CnAlgo<ALGO> props;
|
||||
constexpr Algorithm::Id BASE = props.base();
|
||||
|
||||
if (props.isBase1() && size < 43) {
|
||||
if (BASE == Algorithm::CN_1 && size < 43) {
|
||||
memset(output, 0, 32);
|
||||
return;
|
||||
}
|
||||
@@ -1150,12 +1163,12 @@ static NOINLINE void cryptonight_single_hash_gr_sse41(const uint8_t* __restrict_
|
||||
VARIANT1_INIT(0);
|
||||
ctx[0]->tweak1_2 = tweak1_2_0;
|
||||
ctx[0]->tweak1_table = tweak1_table;
|
||||
if (props.isGR0()) cn_gr0_single_mainloop_asm(ctx);
|
||||
if (props.isGR1()) cn_gr1_single_mainloop_asm(ctx);
|
||||
if (props.isGR2()) cn_gr2_single_mainloop_asm(ctx);
|
||||
if (props.isGR3()) cn_gr3_single_mainloop_asm(ctx);
|
||||
if (props.isGR4()) cn_gr4_single_mainloop_asm(ctx);
|
||||
if (props.isGR5()) cn_gr5_single_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_0) cn_gr0_single_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_1) cn_gr1_single_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_2) cn_gr2_single_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_3) cn_gr3_single_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_4) cn_gr4_single_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_5) cn_gr5_single_mainloop_asm(ctx);
|
||||
|
||||
cn_implode_scratchpad<ALGO, false, 0>(ctx[0]);
|
||||
keccakf(reinterpret_cast<uint64_t*>(ctx[0]->state), 24);
|
||||
@@ -1167,8 +1180,9 @@ template<Algorithm::Id ALGO>
|
||||
static NOINLINE void cryptonight_double_hash_gr_sse41(const uint8_t *__restrict__ input, size_t size, uint8_t *__restrict__ output, cryptonight_ctx **__restrict__ ctx, uint64_t height)
|
||||
{
|
||||
constexpr CnAlgo<ALGO> props;
|
||||
constexpr Algorithm::Id BASE = props.base();
|
||||
|
||||
if (props.isBase1() && size < 43) {
|
||||
if (BASE == Algorithm::CN_1 && size < 43) {
|
||||
memset(output, 0, 64);
|
||||
return;
|
||||
}
|
||||
@@ -1182,7 +1196,7 @@ static NOINLINE void cryptonight_double_hash_gr_sse41(const uint8_t *__restrict_
|
||||
}
|
||||
|
||||
# ifdef XMRIG_VAES
|
||||
if (cn_vaes_enabled) {
|
||||
if (!props.isHeavy() && cn_vaes_enabled) {
|
||||
cn_explode_scratchpad_vaes_double(ctx[0], ctx[1], props.memory(), props.half_mem());
|
||||
}
|
||||
else
|
||||
@@ -1200,15 +1214,15 @@ static NOINLINE void cryptonight_double_hash_gr_sse41(const uint8_t *__restrict_
|
||||
|
||||
ctx[0]->tweak1_table = tweak1_table;
|
||||
|
||||
if (props.isGR0()) cn_gr0_double_mainloop_asm(ctx);
|
||||
if (props.isGR1()) cn_gr1_double_mainloop_asm(ctx);
|
||||
if (props.isGR2()) cn_gr2_double_mainloop_asm(ctx);
|
||||
if (props.isGR3()) cn_gr3_double_mainloop_asm(ctx);
|
||||
if (props.isGR4()) cn_gr4_double_mainloop_asm(ctx);
|
||||
if (props.isGR5()) cn_gr5_double_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_0) cn_gr0_double_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_1) cn_gr1_double_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_2) cn_gr2_double_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_3) cn_gr3_double_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_4) cn_gr4_double_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_5) cn_gr5_double_mainloop_asm(ctx);
|
||||
|
||||
# ifdef XMRIG_VAES
|
||||
if (cn_vaes_enabled) {
|
||||
if (!props.isHeavy() && cn_vaes_enabled) {
|
||||
cn_implode_scratchpad_vaes_double(ctx[0], ctx[1], props.memory(), props.half_mem());
|
||||
}
|
||||
else
|
||||
@@ -1253,8 +1267,15 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
|
||||
|
||||
constexpr CnAlgo<ALGO> props;
|
||||
constexpr size_t MASK = props.mask();
|
||||
constexpr Algorithm::Id BASE = props.base();
|
||||
|
||||
if (props.isBase1() && size < 43) {
|
||||
# ifdef XMRIG_ALGO_CN_HEAVY
|
||||
constexpr bool IS_CN_HEAVY_TUBE = ALGO == Algorithm::CN_HEAVY_TUBE;
|
||||
# else
|
||||
constexpr bool IS_CN_HEAVY_TUBE = false;
|
||||
# endif
|
||||
|
||||
if (BASE == Algorithm::CN_1 && size < 43) {
|
||||
memset(output, 0, 64);
|
||||
return;
|
||||
}
|
||||
@@ -1302,7 +1323,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
|
||||
__m128i bx11 = _mm_set_epi64x(h1[9] ^ h1[11], h1[8] ^ h1[10]);
|
||||
|
||||
__m128 conc_var0, conc_var1;
|
||||
if (props.isCCX()) {
|
||||
if (ALGO == Algorithm::CN_CCX) {
|
||||
conc_var0 = _mm_setzero_ps();
|
||||
conc_var1 = _mm_setzero_ps();
|
||||
RESTORE_ROUNDING_MODE();
|
||||
@@ -1313,10 +1334,10 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
|
||||
|
||||
for (size_t i = 0; i < props.iterations(); i++) {
|
||||
__m128i cx0, cx1;
|
||||
if (props.isHeavyTube() || !SOFT_AES) {
|
||||
if (IS_CN_HEAVY_TUBE || !SOFT_AES) {
|
||||
cx0 = _mm_load_si128(reinterpret_cast<const __m128i *>(&l0[idx0 & MASK]));
|
||||
cx1 = _mm_load_si128(reinterpret_cast<const __m128i *>(&l1[idx1 & MASK]));
|
||||
if (props.isCCX()) {
|
||||
if (ALGO == Algorithm::CN_CCX) {
|
||||
cryptonight_conceal_tweak(cx0, conc_var0);
|
||||
cryptonight_conceal_tweak(cx1, conc_var1);
|
||||
}
|
||||
@@ -1324,12 +1345,12 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
|
||||
|
||||
const __m128i ax0 = _mm_set_epi64x(ah0, al0);
|
||||
const __m128i ax1 = _mm_set_epi64x(ah1, al1);
|
||||
if (props.isHeavyTube()) {
|
||||
if (IS_CN_HEAVY_TUBE) {
|
||||
cx0 = aes_round_tweak_div(cx0, ax0);
|
||||
cx1 = aes_round_tweak_div(cx1, ax1);
|
||||
}
|
||||
else if (SOFT_AES) {
|
||||
if (props.isCCX()) {
|
||||
if (ALGO == Algorithm::CN_CCX) {
|
||||
cx0 = _mm_load_si128(reinterpret_cast<const __m128i*>(&l0[idx0 & MASK]));
|
||||
cx1 = _mm_load_si128(reinterpret_cast<const __m128i*>(&l1[idx1 & MASK]));
|
||||
cryptonight_conceal_tweak(cx0, conc_var0);
|
||||
@@ -1347,7 +1368,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
|
||||
cx1 = _mm_aesenc_si128(cx1, ax1);
|
||||
}
|
||||
|
||||
if (props.isBase1() || props.isBase2()) {
|
||||
if (BASE == Algorithm::CN_1 || BASE == Algorithm::CN_2) {
|
||||
cryptonight_monero_tweak<ALGO>((uint64_t*)&l0[idx0 & MASK], l0, idx0 & MASK, ax0, bx00, bx01, cx0);
|
||||
cryptonight_monero_tweak<ALGO>((uint64_t*)&l1[idx1 & MASK], l1, idx1 & MASK, ax1, bx10, bx11, cx1);
|
||||
} else {
|
||||
@@ -1362,11 +1383,13 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
|
||||
cl = ((uint64_t*) &l0[idx0 & MASK])[0];
|
||||
ch = ((uint64_t*) &l0[idx0 & MASK])[1];
|
||||
|
||||
if (props.isBase2()) {
|
||||
if (BASE == Algorithm::CN_2) {
|
||||
if (props.isR()) {
|
||||
VARIANT4_RANDOM_MATH(0, al0, ah0, cl, bx00, bx01);
|
||||
al0 ^= r0[2] | ((uint64_t)(r0[3]) << 32);
|
||||
ah0 ^= r0[0] | ((uint64_t)(r0[1]) << 32);
|
||||
if (ALGO == Algorithm::CN_R) {
|
||||
al0 ^= r0[2] | ((uint64_t)(r0[3]) << 32);
|
||||
ah0 ^= r0[0] | ((uint64_t)(r0[1]) << 32);
|
||||
}
|
||||
} else {
|
||||
VARIANT2_INTEGER_MATH(0, cl, cx0);
|
||||
}
|
||||
@@ -1374,11 +1397,11 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
|
||||
|
||||
lo = __umul128(idx0, cl, &hi);
|
||||
|
||||
if (props.isBase2()) {
|
||||
if (props.isR()) {
|
||||
if (BASE == Algorithm::CN_2) {
|
||||
if (ALGO == Algorithm::CN_R) {
|
||||
VARIANT2_SHUFFLE(l0, idx0 & MASK, ax0, bx00, bx01, cx0, 0);
|
||||
} else {
|
||||
VARIANT2_SHUFFLE2(l0, idx0 & MASK, ax0, bx00, bx01, hi, lo, ((props.isRWZ() || props.isUPX2()) ? 1 : 0));
|
||||
VARIANT2_SHUFFLE2(l0, idx0 & MASK, ax0, bx00, bx01, hi, lo, (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? 1 : 0));
|
||||
}
|
||||
}
|
||||
|
||||
@@ -1387,9 +1410,9 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
|
||||
|
||||
((uint64_t*)&l0[idx0 & MASK])[0] = al0;
|
||||
|
||||
if (props.isHeavyTube() || props.isRTO()) {
|
||||
if (IS_CN_HEAVY_TUBE || ALGO == Algorithm::CN_RTO) {
|
||||
((uint64_t*) &l0[idx0 & MASK])[1] = ah0 ^ tweak1_2_0 ^ al0;
|
||||
} else if (props.isBase1()) {
|
||||
} else if (BASE == Algorithm::CN_1) {
|
||||
((uint64_t*) &l0[idx0 & MASK])[1] = ah0 ^ tweak1_2_0;
|
||||
} else {
|
||||
((uint64_t*) &l0[idx0 & MASK])[1] = ah0;
|
||||
@@ -1407,7 +1430,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
|
||||
|
||||
((int64_t*)&l0[idx0 & MASK])[0] = n ^ q;
|
||||
|
||||
if (props.isHeavyXHV()) {
|
||||
if (ALGO == Algorithm::CN_HEAVY_XHV) {
|
||||
d = ~d;
|
||||
}
|
||||
|
||||
@@ -1418,11 +1441,13 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
|
||||
cl = ((uint64_t*) &l1[idx1 & MASK])[0];
|
||||
ch = ((uint64_t*) &l1[idx1 & MASK])[1];
|
||||
|
||||
if (props.isBase2()) {
|
||||
if (BASE == Algorithm::CN_2) {
|
||||
if (props.isR()) {
|
||||
VARIANT4_RANDOM_MATH(1, al1, ah1, cl, bx10, bx11);
|
||||
al1 ^= r1[2] | ((uint64_t)(r1[3]) << 32);
|
||||
ah1 ^= r1[0] | ((uint64_t)(r1[1]) << 32);
|
||||
if (ALGO == Algorithm::CN_R) {
|
||||
al1 ^= r1[2] | ((uint64_t)(r1[3]) << 32);
|
||||
ah1 ^= r1[0] | ((uint64_t)(r1[1]) << 32);
|
||||
}
|
||||
} else {
|
||||
VARIANT2_INTEGER_MATH(1, cl, cx1);
|
||||
}
|
||||
@@ -1430,11 +1455,11 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
|
||||
|
||||
lo = __umul128(idx1, cl, &hi);
|
||||
|
||||
if (props.isBase2()) {
|
||||
if (props.isR()) {
|
||||
if (BASE == Algorithm::CN_2) {
|
||||
if (ALGO == Algorithm::CN_R) {
|
||||
VARIANT2_SHUFFLE(l1, idx1 & MASK, ax1, bx10, bx11, cx1, 0);
|
||||
} else {
|
||||
VARIANT2_SHUFFLE2(l1, idx1 & MASK, ax1, bx10, bx11, hi, lo, ((props.isRWZ() || props.isUPX2()) ? 1 : 0));
|
||||
VARIANT2_SHUFFLE2(l1, idx1 & MASK, ax1, bx10, bx11, hi, lo, (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? 1 : 0));
|
||||
}
|
||||
}
|
||||
|
||||
@@ -1443,9 +1468,9 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
|
||||
|
||||
((uint64_t*)&l1[idx1 & MASK])[0] = al1;
|
||||
|
||||
if (props.isHeavyTube() || props.isRTO()) {
|
||||
if (IS_CN_HEAVY_TUBE || ALGO == Algorithm::CN_RTO) {
|
||||
((uint64_t*)&l1[idx1 & MASK])[1] = ah1 ^ tweak1_2_1 ^ al1;
|
||||
} else if (props.isBase1()) {
|
||||
} else if (BASE == Algorithm::CN_1) {
|
||||
((uint64_t*)&l1[idx1 & MASK])[1] = ah1 ^ tweak1_2_1;
|
||||
} else {
|
||||
((uint64_t*)&l1[idx1 & MASK])[1] = ah1;
|
||||
@@ -1463,7 +1488,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
|
||||
|
||||
((int64_t*)&l1[idx1 & MASK])[0] = n ^ q;
|
||||
|
||||
if (props.isHeavyXHV()) {
|
||||
if (ALGO == Algorithm::CN_HEAVY_XHV) {
|
||||
d = ~d;
|
||||
}
|
||||
|
||||
@@ -1471,7 +1496,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
|
||||
}
|
||||
# endif
|
||||
|
||||
if (props.isBase2()) {
|
||||
if (BASE == Algorithm::CN_2) {
|
||||
bx01 = bx00;
|
||||
bx11 = bx10;
|
||||
}
|
||||
@@ -1504,8 +1529,9 @@ template<Algorithm::Id ALGO>
|
||||
static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__ input, size_t size, uint8_t* __restrict__ output, cryptonight_ctx** __restrict__ ctx, uint64_t height)
|
||||
{
|
||||
constexpr CnAlgo<ALGO> props;
|
||||
constexpr Algorithm::Id BASE = props.base();
|
||||
|
||||
if (props.isBase1() && size < 43) {
|
||||
if (BASE == Algorithm::CN_1 && size < 43) {
|
||||
memset(output, 0, 32 * 4);
|
||||
return;
|
||||
}
|
||||
@@ -1523,7 +1549,7 @@ static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__
|
||||
}
|
||||
|
||||
# ifdef XMRIG_VAES
|
||||
if (cn_vaes_enabled) {
|
||||
if (!props.isHeavy() && cn_vaes_enabled) {
|
||||
cn_explode_scratchpad_vaes_double(ctx[0], ctx[1], props.memory(), props.half_mem());
|
||||
cn_explode_scratchpad_vaes_double(ctx[2], ctx[3], props.memory(), props.half_mem());
|
||||
}
|
||||
@@ -1543,15 +1569,15 @@ static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__
|
||||
|
||||
ctx[0]->tweak1_table = tweak1_table;
|
||||
|
||||
if (props.isGR0()) cn_gr0_quad_mainloop_asm(ctx);
|
||||
if (props.isGR1()) cn_gr1_quad_mainloop_asm(ctx);
|
||||
if (props.isGR2()) cn_gr2_quad_mainloop_asm(ctx);
|
||||
if (props.isGR3()) cn_gr3_quad_mainloop_asm(ctx);
|
||||
if (props.isGR4()) cn_gr4_quad_mainloop_asm(ctx);
|
||||
if (props.isGR5()) cn_gr5_quad_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_0) cn_gr0_quad_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_1) cn_gr1_quad_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_2) cn_gr2_quad_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_3) cn_gr3_quad_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_4) cn_gr4_quad_mainloop_asm(ctx);
|
||||
if (ALGO == Algorithm::CN_GR_5) cn_gr5_quad_mainloop_asm(ctx);
|
||||
|
||||
# ifdef XMRIG_VAES
|
||||
if (cn_vaes_enabled) {
|
||||
if (!props.isHeavy() && cn_vaes_enabled) {
|
||||
cn_implode_scratchpad_vaes_double(ctx[0], ctx[1], props.memory(), props.half_mem());
|
||||
cn_implode_scratchpad_vaes_double(ctx[2], ctx[3], props.memory(), props.half_mem());
|
||||
}
|
||||
@@ -1580,14 +1606,14 @@ static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__
|
||||
#define CN_STEP1(a, b0, b1, c, l, ptr, idx, conc_var) \
|
||||
ptr = reinterpret_cast<__m128i*>(&l[idx & MASK]); \
|
||||
c = _mm_load_si128(ptr); \
|
||||
if (props.isCCX()) { \
|
||||
if (ALGO == Algorithm::CN_CCX) { \
|
||||
cryptonight_conceal_tweak(c, conc_var); \
|
||||
}
|
||||
|
||||
|
||||
|
||||
#define CN_STEP2(a, b0, b1, c, l, ptr, idx) \
|
||||
if (props.isHeavyTube()) { \
|
||||
if (IS_CN_HEAVY_TUBE) { \
|
||||
c = aes_round_tweak_div(c, a); \
|
||||
} \
|
||||
else if (SOFT_AES) { \
|
||||
@@ -1596,7 +1622,7 @@ static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__
|
||||
c = _mm_aesenc_si128(c, a); \
|
||||
} \
|
||||
\
|
||||
if (props.isBase1() || props.isBase2()) { \
|
||||
if (BASE == Algorithm::CN_1 || BASE == Algorithm::CN_2) { \
|
||||
cryptonight_monero_tweak<ALGO>((uint64_t*)ptr, l, idx & MASK, a, b0, b1, c); \
|
||||
} else { \
|
||||
_mm_store_si128(ptr, _mm_xor_si128(b0, c)); \
|
||||
@@ -1612,34 +1638,36 @@ static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__
|
||||
|
||||
#define CN_STEP4(part, a, b0, b1, c, l, mc, ptr, idx) \
|
||||
uint64_t al##part, ah##part; \
|
||||
if (props.isBase2()) { \
|
||||
if (BASE == Algorithm::CN_2) { \
|
||||
if (props.isR()) { \
|
||||
al##part = _mm_cvtsi128_si64(a); \
|
||||
ah##part = _mm_cvtsi128_si64(_mm_srli_si128(a, 8)); \
|
||||
VARIANT4_RANDOM_MATH(part, al##part, ah##part, cl##part, b0, b1); \
|
||||
al##part ^= r##part[2] | ((uint64_t)(r##part[3]) << 32); \
|
||||
ah##part ^= r##part[0] | ((uint64_t)(r##part[1]) << 32); \
|
||||
if (ALGO == Algorithm::CN_R) { \
|
||||
al##part ^= r##part[2] | ((uint64_t)(r##part[3]) << 32); \
|
||||
ah##part ^= r##part[0] | ((uint64_t)(r##part[1]) << 32); \
|
||||
} \
|
||||
} else { \
|
||||
VARIANT2_INTEGER_MATH(part, cl##part, c); \
|
||||
} \
|
||||
} \
|
||||
lo = __umul128(idx, cl##part, &hi); \
|
||||
if (props.isBase2()) { \
|
||||
if (props.isR()) { \
|
||||
if (BASE == Algorithm::CN_2) { \
|
||||
if (ALGO == Algorithm::CN_R) { \
|
||||
VARIANT2_SHUFFLE(l, idx & MASK, a, b0, b1, c, 0); \
|
||||
} else { \
|
||||
VARIANT2_SHUFFLE2(l, idx & MASK, a, b0, b1, hi, lo, ((props.isRWZ() || props.isUPX2()) ? 1 : 0)); \
|
||||
VARIANT2_SHUFFLE2(l, idx & MASK, a, b0, b1, hi, lo, (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? 1 : 0)); \
|
||||
} \
|
||||
} \
|
||||
if (props.isR()) { \
|
||||
if (ALGO == Algorithm::CN_R) { \
|
||||
a = _mm_set_epi64x(ah##part, al##part); \
|
||||
} \
|
||||
a = _mm_add_epi64(a, _mm_set_epi64x(lo, hi)); \
|
||||
\
|
||||
if (props.isBase1()) { \
|
||||
if (BASE == Algorithm::CN_1) { \
|
||||
_mm_store_si128(ptr, _mm_xor_si128(a, mc)); \
|
||||
\
|
||||
if (props.isHeavyTube() || props.isRTO()) { \
|
||||
if (IS_CN_HEAVY_TUBE || ALGO == Algorithm::CN_RTO) { \
|
||||
((uint64_t*)ptr)[1] ^= ((uint64_t*)ptr)[0]; \
|
||||
} \
|
||||
} else { \
|
||||
@@ -1653,13 +1681,13 @@ static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__
|
||||
int32_t d = ((int32_t*)&l[idx & MASK])[2]; \
|
||||
int64_t q = n / (d | 0x5); \
|
||||
((int64_t*)&l[idx & MASK])[0] = n ^ q; \
|
||||
if (props.isHeavyXHV()) { \
|
||||
if (IS_CN_HEAVY_XHV) { \
|
||||
d = ~d; \
|
||||
} \
|
||||
\
|
||||
idx = d ^ q; \
|
||||
} \
|
||||
if (props.isBase2()) { \
|
||||
if (BASE == Algorithm::CN_2) { \
|
||||
b1 = b0; \
|
||||
} \
|
||||
b0 = c;
|
||||
@@ -1669,11 +1697,11 @@ static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__
|
||||
__m128i mc##n; \
|
||||
__m128i division_result_xmm_##n; \
|
||||
__m128i sqrt_result_xmm_##n; \
|
||||
if (props.isBase1()) { \
|
||||
if (BASE == Algorithm::CN_1) { \
|
||||
mc##n = _mm_set_epi64x(*reinterpret_cast<const uint64_t*>(input + n * size + 35) ^ \
|
||||
*(reinterpret_cast<const uint64_t*>((ctx)->state) + 24), 0); \
|
||||
} \
|
||||
if (props.isBase2()) { \
|
||||
if (BASE == Algorithm::CN_2) { \
|
||||
division_result_xmm_##n = _mm_cvtsi64_si128(h##n[12]); \
|
||||
sqrt_result_xmm_##n = _mm_cvtsi64_si128(h##n[13]); \
|
||||
} \
|
||||
@@ -1682,7 +1710,7 @@ static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__
|
||||
__m128i bx##n##1 = _mm_set_epi64x(h##n[9] ^ h##n[11], h##n[8] ^ h##n[10]); \
|
||||
__m128i cx##n = _mm_setzero_si128(); \
|
||||
__m128 conc_var##n; \
|
||||
if (props.isCCX()) { \
|
||||
if (ALGO == Algorithm::CN_CCX) { \
|
||||
conc_var##n = _mm_setzero_ps(); \
|
||||
} \
|
||||
VARIANT4_RANDOM_MATH_INIT(n);
|
||||
@@ -1693,8 +1721,17 @@ inline void cryptonight_triple_hash(const uint8_t *__restrict__ input, size_t si
|
||||
{
|
||||
constexpr CnAlgo<ALGO> props;
|
||||
constexpr size_t MASK = props.mask();
|
||||
constexpr Algorithm::Id BASE = props.base();
|
||||
|
||||
if (props.isBase1() && size < 43) {
|
||||
# ifdef XMRIG_ALGO_CN_HEAVY
|
||||
constexpr bool IS_CN_HEAVY_TUBE = ALGO == Algorithm::CN_HEAVY_TUBE;
|
||||
constexpr bool IS_CN_HEAVY_XHV = ALGO == Algorithm::CN_HEAVY_XHV;
|
||||
# else
|
||||
constexpr bool IS_CN_HEAVY_TUBE = false;
|
||||
constexpr bool IS_CN_HEAVY_XHV = false;
|
||||
# endif
|
||||
|
||||
if (BASE == Algorithm::CN_1 && size < 43) {
|
||||
memset(output, 0, 32 * 3);
|
||||
return;
|
||||
}
|
||||
@@ -1718,7 +1755,7 @@ inline void cryptonight_triple_hash(const uint8_t *__restrict__ input, size_t si
|
||||
CONST_INIT(ctx[1], 1);
|
||||
CONST_INIT(ctx[2], 2);
|
||||
VARIANT2_SET_ROUNDING_MODE();
|
||||
if (props.isCCX()) {
|
||||
if (ALGO == Algorithm::CN_CCX) {
|
||||
RESTORE_ROUNDING_MODE();
|
||||
}
|
||||
|
||||
@@ -1782,8 +1819,17 @@ inline void cryptonight_quad_hash(const uint8_t *__restrict__ input, size_t size
|
||||
|
||||
constexpr CnAlgo<ALGO> props;
|
||||
constexpr size_t MASK = props.mask();
|
||||
constexpr Algorithm::Id BASE = props.base();
|
||||
|
||||
if (props.isBase1() && size < 43) {
|
||||
# ifdef XMRIG_ALGO_CN_HEAVY
|
||||
constexpr bool IS_CN_HEAVY_TUBE = ALGO == Algorithm::CN_HEAVY_TUBE;
|
||||
constexpr bool IS_CN_HEAVY_XHV = ALGO == Algorithm::CN_HEAVY_XHV;
|
||||
# else
|
||||
constexpr bool IS_CN_HEAVY_TUBE = false;
|
||||
constexpr bool IS_CN_HEAVY_XHV = false;
|
||||
# endif
|
||||
|
||||
if (BASE == Algorithm::CN_1 && size < 43) {
|
||||
memset(output, 0, 32 * 4);
|
||||
return;
|
||||
}
|
||||
@@ -1823,7 +1869,7 @@ inline void cryptonight_quad_hash(const uint8_t *__restrict__ input, size_t size
|
||||
CONST_INIT(ctx[2], 2);
|
||||
CONST_INIT(ctx[3], 3);
|
||||
VARIANT2_SET_ROUNDING_MODE();
|
||||
if (props.isCCX()) {
|
||||
if (ALGO == Algorithm::CN_CCX) {
|
||||
RESTORE_ROUNDING_MODE();
|
||||
}
|
||||
|
||||
@@ -1884,8 +1930,17 @@ inline void cryptonight_penta_hash(const uint8_t *__restrict__ input, size_t siz
|
||||
{
|
||||
constexpr CnAlgo<ALGO> props;
|
||||
constexpr size_t MASK = props.mask();
|
||||
constexpr Algorithm::Id BASE = props.base();
|
||||
|
||||
if (props.isBase1() && size < 43) {
|
||||
# ifdef XMRIG_ALGO_CN_HEAVY
|
||||
constexpr bool IS_CN_HEAVY_TUBE = ALGO == Algorithm::CN_HEAVY_TUBE;
|
||||
constexpr bool IS_CN_HEAVY_XHV = ALGO == Algorithm::CN_HEAVY_XHV;
|
||||
# else
|
||||
constexpr bool IS_CN_HEAVY_TUBE = false;
|
||||
constexpr bool IS_CN_HEAVY_XHV = false;
|
||||
# endif
|
||||
|
||||
if (BASE == Algorithm::CN_1 && size < 43) {
|
||||
memset(output, 0, 32 * 5);
|
||||
return;
|
||||
}
|
||||
@@ -1915,7 +1970,7 @@ inline void cryptonight_penta_hash(const uint8_t *__restrict__ input, size_t siz
|
||||
CONST_INIT(ctx[3], 3);
|
||||
CONST_INIT(ctx[4], 4);
|
||||
VARIANT2_SET_ROUNDING_MODE();
|
||||
if (props.isCCX()) {
|
||||
if (ALGO == Algorithm::CN_CCX) {
|
||||
RESTORE_ROUNDING_MODE();
|
||||
}
|
||||
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@@ -1,6 +1,6 @@
|
||||
/* XMRig
|
||||
* Copyright (c) 2018-2023 SChernykh <https://github.com/SChernykh>
|
||||
* Copyright (c) 2016-2023 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||
* Copyright (c) 2018-2024 SChernykh <https://github.com/SChernykh>
|
||||
* Copyright (c) 2016-2024 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||
*
|
||||
* This program is free software: you can redistribute it and/or modify
|
||||
* it under the terms of the GNU General Public License as published by
|
||||
@@ -22,14 +22,14 @@
|
||||
#define APP_ID "xmrig"
|
||||
#define APP_NAME "XMRig"
|
||||
#define APP_DESC "XMRig miner"
|
||||
#define APP_VERSION "6.20.1-dev"
|
||||
#define APP_VERSION "6.21.1"
|
||||
#define APP_DOMAIN "xmrig.com"
|
||||
#define APP_SITE "www.xmrig.com"
|
||||
#define APP_COPYRIGHT "Copyright (C) 2016-2023 xmrig.com"
|
||||
#define APP_COPYRIGHT "Copyright (C) 2016-2024 xmrig.com"
|
||||
#define APP_KIND "miner"
|
||||
|
||||
#define APP_VER_MAJOR 6
|
||||
#define APP_VER_MINOR 20
|
||||
#define APP_VER_MINOR 21
|
||||
#define APP_VER_PATCH 1
|
||||
|
||||
#ifdef _MSC_VER
|
||||
|
||||
Reference in New Issue
Block a user