Skip to content

Commit

Permalink
Merged xmrig v6.21.1
Browse files Browse the repository at this point in the history
  • Loading branch information
@MoneroOcean
MoneroOcean committed Feb 25, 2024
2 parents 06c383e + a5aa2c9 commit 923dbdf
Show file tree
Hide file tree
Showing 16 changed files with 2,855 additions and 2,490 deletions.
5 changes: 5 additions & 0 deletions CHANGELOG.md
View file
@@ -1,3 +1,8 @@
# 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".
Expand Down
87 changes: 87 additions & 0 deletions src/backend/opencl/cl/rx/blake2b.cl
View file
Expand Up @@ -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
Expand Down
4,966 changes: 2,518 additions & 2,448 deletions src/backend/opencl/cl/rx/randomx_cl.h
View file

Large diffs are not rendered by default.

58 changes: 58 additions & 0 deletions src/backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.cpp
View file
@@ -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, &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);
}
50 changes: 50 additions & 0 deletions src/backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.h
View file
@@ -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 */
2 changes: 2 additions & 0 deletions src/backend/opencl/opencl.cmake
View file
Expand Up @@ -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
Expand All @@ -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
Expand Down
29 changes: 27 additions & 2 deletions src/backend/opencl/runners/OclRxBaseRunner.cpp
View file
Expand Up @@ -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"
Expand Down Expand Up @@ -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;
Expand All @@ -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);
Expand Down Expand Up @@ -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());
}

Expand Down Expand Up @@ -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");

Expand Down
32 changes: 18 additions & 14 deletions src/backend/opencl/runners/OclRxBaseRunner.h
View file
Expand Up @@ -35,6 +35,7 @@ namespace xmrig {

class Blake2bHashRegistersKernel;
class Blake2bInitialHashKernel;
class Blake2bInitialHashDoubleKernel;
class FillAesKernel;
class FindSharesKernel;
class HashAesKernel;
Expand All @@ -58,21 +59,24 @@ class OclRxBaseRunner : public OclBaseRunner
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;
};


Expand Down

0 comments on commit 923dbdf

Please sign in to comment.