Merge pull request #3528 from SChernykh/dev
Added rx/yada OpenCL support
This commit is contained in:
xmrig
GitHub
commit
c5f98fc5c7
@ -170,7 +170,7 @@ void xmrig::OclWorker::start()
|
||||
const uint64_t t = Chrono::steadyMSecs();
|
||||
|
||||
try {
|
||||
m_runner->run(readUnaligned(m_job.nonce()), results);
|
||||
m_runner->run(readUnaligned(m_job.nonce()), m_job.nonceOffset(), results);
|
||||
}
|
||||
catch (std::exception &ex) {
|
||||
printError(id(), ex.what());
|
||||
|
||||
@ -23,6 +23,7 @@
|
||||
#define ALGO_RX_ARQMA 0x72121061
|
||||
#define ALGO_RX_SFX 0x72151273
|
||||
#define ALGO_RX_GRAFT 0x72151267
|
||||
#define ALGO_RX_YADA 0x72151279
|
||||
#define ALGO_AR2_CHUKWA 0x61130000
|
||||
#define ALGO_AR2_CHUKWA_V2 0x61140000
|
||||
#define ALGO_AR2_WRKZ 0x61120000
|
||||
|
||||
@ -2,7 +2,7 @@
|
||||
|
||||
namespace xmrig {
|
||||
|
||||
static const char cryptonight_cl[61415] = {
|
||||
static const char cryptonight_cl[61447] = {
|
||||
0x23,0x69,0x66,0x64,0x65,0x66,0x20,0x53,0x54,0x41,0x54,0x49,0x43,0x0a,0x23,0x75,0x6e,0x64,0x65,0x66,0x20,0x53,0x54,0x41,0x54,0x49,0x43,0x0a,0x23,0x65,0x6e,0x64,
|
||||
0x69,0x66,0x0a,0x23,0x69,0x66,0x64,0x65,0x66,0x20,0x63,0x6c,0x5f,0x61,0x6d,0x64,0x5f,0x6d,0x65,0x64,0x69,0x61,0x5f,0x6f,0x70,0x73,0x0a,0x23,0x64,0x65,0x66,0x69,
|
||||
0x6e,0x65,0x20,0x53,0x54,0x41,0x54,0x49,0x43,0x20,0x73,0x74,0x61,0x74,0x69,0x63,0x0a,0x23,0x65,0x6c,0x73,0x65,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x53,
|
||||
@ -35,7 +35,8 @@ static const char cryptonight_cl[61415] = {
|
||||
0x34,0x31,0x31,0x37,0x37,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x41,0x52,0x51,0x4d,0x41,0x20,0x30,0x78,0x37,0x32,
|
||||
0x31,0x32,0x31,0x30,0x36,0x31,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x53,0x46,0x58,0x20,0x30,0x78,0x37,0x32,0x31,
|
||||
0x35,0x31,0x32,0x37,0x33,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x47,0x52,0x41,0x46,0x54,0x20,0x30,0x78,0x37,0x32,
|
||||
0x31,0x35,0x31,0x32,0x36,0x37,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,0x41,0x20,0x30,
|
||||
0x31,0x35,0x31,0x32,0x36,0x37,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x59,0x41,0x44,0x41,0x20,0x30,0x78,0x37,0x32,
|
||||
0x31,0x35,0x31,0x32,0x37,0x39,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,0x41,0x20,0x30,
|
||||
0x78,0x36,0x31,0x31,0x33,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,
|
||||
0x41,0x5f,0x56,0x32,0x20,0x30,0x78,0x36,0x31,0x31,0x34,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,
|
||||
0x5f,0x57,0x52,0x4b,0x5a,0x20,0x30,0x78,0x36,0x31,0x31,0x32,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x4b,0x41,
|
||||
|
||||
@ -225,6 +225,110 @@ __kernel void blake2b_initial_hash_double(__global void *out, __global const voi
|
||||
t[7] = hash[7];
|
||||
}
|
||||
|
||||
void blake2b_512_process_big_block(ulong *out, __global const ulong* in, uint in_len, uint out_len, uint nonce, uint nonce_offset)
|
||||
{
|
||||
ulong h[8] = { iv0 ^ (0x01010000u | out_len), iv1, iv2, iv3, iv4, iv5, iv6, iv7 };
|
||||
|
||||
for (uint t = 128; t < in_len; t += 128, in += 16) {
|
||||
ulong m[16] = { in[0], in[1], in[2], in[3], in[4], in[5], in[6], in[7], in[8], in[9], in[10], in[11], in[12], in[13], in[14], in[15] };
|
||||
|
||||
const uint k0 = (nonce_offset + 0) - (t - 128);
|
||||
const uint k1 = (nonce_offset + 1) - (t - 128);
|
||||
const uint k2 = (nonce_offset + 2) - (t - 128);
|
||||
const uint k3 = (nonce_offset + 3) - (t - 128);
|
||||
|
||||
if (k0 < 128) m[k0 / 8] |= (ulong)((nonce >> 0) & 255) << ((k0 % 8) * 8);
|
||||
if (k1 < 128) m[k1 / 8] |= (ulong)((nonce >> 8) & 255) << ((k1 % 8) * 8);
|
||||
if (k2 < 128) m[k2 / 8] |= (ulong)((nonce >> 16) & 255) << ((k2 % 8) * 8);
|
||||
if (k3 < 128) m[k3 / 8] |= (ulong)((nonce >> 24) & 255) << ((k3 % 8) * 8);
|
||||
|
||||
ulong v[16] = { h[0], h[1], h[2], h[3], h[4], h[5], h[6], h[7], iv0, iv1, iv2, iv3, iv4 ^ t, iv5, iv6, iv7 };
|
||||
|
||||
BLAKE2B_ROUNDS();
|
||||
|
||||
h[0] ^= v[0] ^ v[ 8];
|
||||
h[1] ^= v[1] ^ v[ 9];
|
||||
h[2] ^= v[2] ^ v[10];
|
||||
h[3] ^= v[3] ^ v[11];
|
||||
h[4] ^= v[4] ^ v[12];
|
||||
h[5] ^= v[5] ^ v[13];
|
||||
h[6] ^= v[6] ^ v[14];
|
||||
h[7] ^= v[7] ^ v[15];
|
||||
}
|
||||
|
||||
uint k = in_len & 127;
|
||||
if (k == 0) k = 128;
|
||||
|
||||
ulong m[16] = {
|
||||
(k > 0) ? in[ 0] : 0,
|
||||
(k > 8) ? in[ 1] : 0,
|
||||
(k > 16) ? in[ 2] : 0,
|
||||
(k > 24) ? in[ 3] : 0,
|
||||
(k > 32) ? in[ 4] : 0,
|
||||
(k > 40) ? in[ 5] : 0,
|
||||
(k > 48) ? in[ 6] : 0,
|
||||
(k > 56) ? in[ 7] : 0,
|
||||
(k > 64) ? in[ 8] : 0,
|
||||
(k > 72) ? in[ 9] : 0,
|
||||
(k > 80) ? in[10] : 0,
|
||||
(k > 88) ? in[11] : 0,
|
||||
(k > 96) ? in[12] : 0,
|
||||
(k > 104) ? in[13] : 0,
|
||||
(k > 112) ? in[14] : 0,
|
||||
(k > 120) ? in[15] : 0
|
||||
};
|
||||
|
||||
const uint t = in_len - k;
|
||||
|
||||
const uint k0 = nonce_offset + 0 - t;
|
||||
const uint k1 = nonce_offset + 1 - t;
|
||||
const uint k2 = nonce_offset + 2 - t;
|
||||
const uint k3 = nonce_offset + 3 - t;
|
||||
|
||||
if (k0 < k) m[k0 / 8] |= (ulong)((nonce >> 0) & 255) << ((k0 % 8) * 8);
|
||||
if (k1 < k) m[k1 / 8] |= (ulong)((nonce >> 8) & 255) << ((k1 % 8) * 8);
|
||||
if (k2 < k) m[k2 / 8] |= (ulong)((nonce >> 16) & 255) << ((k2 % 8) * 8);
|
||||
if (k3 < k) m[k3 / 8] |= (ulong)((nonce >> 24) & 255) << ((k3 % 8) * 8);
|
||||
|
||||
if (k % 8) {
|
||||
m[k / 8] &= (ulong)(-1) >> (64 - (k % 8) * 8);
|
||||
}
|
||||
|
||||
ulong v[16] = { h[0], h[1], h[2], h[3], h[4], h[5], h[6], h[7], iv0, iv1, iv2, iv3, iv4 ^ in_len, iv5, ~iv6, iv7 };
|
||||
|
||||
BLAKE2B_ROUNDS();
|
||||
|
||||
if (out_len > 0) out[0] = h[0] ^ v[0] ^ v[8];
|
||||
if (out_len > 8) out[1] = h[1] ^ v[1] ^ v[9];
|
||||
if (out_len > 16) out[2] = h[2] ^ v[2] ^ v[10];
|
||||
if (out_len > 24) out[3] = h[3] ^ v[3] ^ v[11];
|
||||
if (out_len > 32) out[4] = h[4] ^ v[4] ^ v[12];
|
||||
if (out_len > 40) out[5] = h[5] ^ v[5] ^ v[13];
|
||||
if (out_len > 48) out[6] = h[6] ^ v[6] ^ v[14];
|
||||
if (out_len > 56) out[7] = h[7] ^ v[7] ^ v[15];
|
||||
}
|
||||
|
||||
__attribute__((reqd_work_group_size(64, 1, 1)))
|
||||
__kernel void blake2b_initial_hash_big(__global void *out, __global const void* blockTemplate, uint blockTemplateSize, uint start_nonce, uint nonce_offset)
|
||||
{
|
||||
const uint global_index = get_global_id(0);
|
||||
|
||||
__global const ulong* p = (__global const ulong*) blockTemplate;
|
||||
|
||||
ulong hash[8];
|
||||
blake2b_512_process_big_block(hash, p, blockTemplateSize, 64, start_nonce + global_index, nonce_offset);
|
||||
|
||||
__global ulong* t = ((__global ulong*) out) + global_index * 8;
|
||||
t[0] = hash[0];
|
||||
t[1] = hash[1];
|
||||
t[2] = hash[2];
|
||||
t[3] = hash[3];
|
||||
t[4] = hash[4];
|
||||
t[5] = hash[5];
|
||||
t[6] = hash[6];
|
||||
t[7] = hash[7];
|
||||
}
|
||||
|
||||
#define in_len 256
|
||||
|
||||
#define out_len 32
|
||||
|
||||
@ -1,6 +1,6 @@
|
||||
#include "../cn/algorithm.cl"
|
||||
|
||||
#if (ALGO == ALGO_RX_0)
|
||||
#if ((ALGO == ALGO_RX_0) || (ALGO == ALGO_RX_YADA))
|
||||
#include "randomx_constants_monero.h"
|
||||
#elif (ALGO == ALGO_RX_WOW)
|
||||
#include "randomx_constants_wow.h"
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@ -64,7 +64,7 @@ public:
|
||||
virtual uint32_t deviceIndex() const = 0;
|
||||
virtual void build() = 0;
|
||||
virtual void init() = 0;
|
||||
virtual void run(uint32_t nonce, uint32_t *hashOutput) = 0;
|
||||
virtual void run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput) = 0;
|
||||
virtual void set(const Job &job, uint8_t *blob) = 0;
|
||||
virtual void jobEarlyNotification(const Job&) = 0;
|
||||
|
||||
|
||||
@ -0,0 +1,59 @@
|
||||
/* XMRig
|
||||
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||
* Copyright 2018-2019 SChernykh <https://github.com/SChernykh>
|
||||
* Copyright 2016-2019 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||
*
|
||||
* This program is free software: you can redistribute it and/or modify
|
||||
* it under the terms of the GNU General Public License as published by
|
||||
* the Free Software Foundation, either version 3 of the License, or
|
||||
* (at your option) any later version.
|
||||
*
|
||||
* This program is distributed in the hope that it will be useful,
|
||||
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
* GNU General Public License for more details.
|
||||
*
|
||||
* You should have received a copy of the GNU General Public License
|
||||
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||
*/
|
||||
|
||||
|
||||
#include "backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.h"
|
||||
#include "backend/opencl/wrappers/OclLib.h"
|
||||
|
||||
|
||||
void xmrig::Blake2bInitialHashBigKernel::enqueue(cl_command_queue queue, size_t threads)
|
||||
{
|
||||
const size_t gthreads = threads;
|
||||
static const size_t lthreads = 64;
|
||||
|
||||
enqueueNDRange(queue, 1, nullptr, >hreads, <hreads);
|
||||
}
|
||||
|
||||
|
||||
// __kernel void blake2b_initial_hash_double(__global void *out, __global const void* blockTemplate, uint blockTemplateSize, uint start_nonce)
|
||||
void xmrig::Blake2bInitialHashBigKernel::setArgs(cl_mem out, cl_mem blockTemplate)
|
||||
{
|
||||
setArg(0, sizeof(cl_mem), &out);
|
||||
setArg(1, sizeof(cl_mem), &blockTemplate);
|
||||
}
|
||||
|
||||
|
||||
void xmrig::Blake2bInitialHashBigKernel::setBlobSize(size_t size)
|
||||
{
|
||||
const uint32_t s = size;
|
||||
|
||||
setArg(2, sizeof(uint32_t), &s);
|
||||
}
|
||||
|
||||
|
||||
void xmrig::Blake2bInitialHashBigKernel::setNonce(uint32_t nonce, uint32_t nonce_offset)
|
||||
{
|
||||
setArg(3, sizeof(uint32_t), &nonce);
|
||||
setArg(4, sizeof(uint32_t), &nonce_offset);
|
||||
}
|
||||
50
src/backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.h
Normal file
50
src/backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.h
Normal 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_BLAKE2BINITIALHASHBIGKERNEL_H
|
||||
#define XMRIG_BLAKE2BINITIALHASHBIGKERNEL_H
|
||||
|
||||
|
||||
#include "backend/opencl/wrappers/OclKernel.h"
|
||||
|
||||
|
||||
namespace xmrig {
|
||||
|
||||
|
||||
class Blake2bInitialHashBigKernel : public OclKernel
|
||||
{
|
||||
public:
|
||||
inline Blake2bInitialHashBigKernel(cl_program program) : OclKernel(program, "blake2b_initial_hash_big") {}
|
||||
|
||||
void enqueue(cl_command_queue queue, size_t threads);
|
||||
void setArgs(cl_mem out, cl_mem blockTemplate);
|
||||
void setBlobSize(size_t size);
|
||||
void setNonce(uint32_t nonce, uint32_t nonce_offset);
|
||||
};
|
||||
|
||||
|
||||
} // namespace xmrig
|
||||
|
||||
|
||||
#endif /* XMRIG_BLAKE2BINITIALHASHBIGKERNEL_H */
|
||||
@ -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/Blake2bInitialHashBigKernel.h
|
||||
src/backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.h
|
||||
src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.h
|
||||
src/backend/opencl/kernels/rx/ExecuteVmKernel.h
|
||||
@ -97,6 +98,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/Blake2bInitialHashBigKernel.cpp
|
||||
src/backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.cpp
|
||||
src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.cpp
|
||||
src/backend/opencl/kernels/rx/ExecuteVmKernel.cpp
|
||||
|
||||
@ -87,7 +87,7 @@ size_t xmrig::OclCnRunner::bufferSize() const
|
||||
}
|
||||
|
||||
|
||||
void xmrig::OclCnRunner::run(uint32_t nonce, uint32_t *hashOutput)
|
||||
void xmrig::OclCnRunner::run(uint32_t nonce, uint32_t /*nonce_offset*/, uint32_t *hashOutput)
|
||||
{
|
||||
static const cl_uint zero = 0;
|
||||
|
||||
|
||||
@ -42,7 +42,7 @@ public:
|
||||
|
||||
protected:
|
||||
size_t bufferSize() const override;
|
||||
void run(uint32_t nonce, uint32_t *hashOutput) override;
|
||||
void run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput) override;
|
||||
void set(const Job &job, uint8_t *blob) override;
|
||||
void build() override;
|
||||
void init() override;
|
||||
|
||||
@ -75,7 +75,7 @@ OclKawPowRunner::~OclKawPowRunner()
|
||||
}
|
||||
|
||||
|
||||
void OclKawPowRunner::run(uint32_t nonce, uint32_t *hashOutput)
|
||||
void OclKawPowRunner::run(uint32_t nonce, uint32_t /*nonce_offset*/, uint32_t *hashOutput)
|
||||
{
|
||||
const size_t local_work_size = m_workGroupSize;
|
||||
const size_t global_work_offset = nonce;
|
||||
|
||||
@ -40,7 +40,7 @@ public:
|
||||
~OclKawPowRunner() override;
|
||||
|
||||
protected:
|
||||
void run(uint32_t nonce, uint32_t *hashOutput) override;
|
||||
void run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput) override;
|
||||
void set(const Job &job, uint8_t *blob) override;
|
||||
void build() override;
|
||||
void init() override;
|
||||
|
||||
@ -26,6 +26,7 @@
|
||||
#include "backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h"
|
||||
#include "backend/opencl/kernels/rx/Blake2bInitialHashKernel.h"
|
||||
#include "backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.h"
|
||||
#include "backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.h"
|
||||
#include "backend/opencl/kernels/rx/FillAesKernel.h"
|
||||
#include "backend/opencl/kernels/rx/FindSharesKernel.h"
|
||||
#include "backend/opencl/kernels/rx/HashAesKernel.h"
|
||||
@ -73,6 +74,7 @@ xmrig::OclRxBaseRunner::~OclRxBaseRunner()
|
||||
delete m_hashAes1Rx4;
|
||||
delete m_blake2b_initial_hash;
|
||||
delete m_blake2b_initial_hash_double;
|
||||
delete m_blake2b_initial_hash_big;
|
||||
delete m_blake2b_hash_registers_32;
|
||||
delete m_blake2b_hash_registers_64;
|
||||
delete m_find_shares;
|
||||
@ -85,7 +87,7 @@ xmrig::OclRxBaseRunner::~OclRxBaseRunner()
|
||||
}
|
||||
|
||||
|
||||
void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
|
||||
void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput)
|
||||
{
|
||||
static const uint32_t zero = 0;
|
||||
|
||||
@ -96,8 +98,7 @@ void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
|
||||
m_blake2b_initial_hash_double->setNonce(nonce);
|
||||
}
|
||||
else {
|
||||
hashOutput[0xFF] = 0;
|
||||
return;
|
||||
m_blake2b_initial_hash_big->setNonce(nonce, nonce_offset);
|
||||
}
|
||||
|
||||
m_find_shares->setNonce(nonce);
|
||||
@ -107,9 +108,12 @@ void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
|
||||
if (m_jobSize <= 128) {
|
||||
m_blake2b_initial_hash->enqueue(m_queue, m_intensity);
|
||||
}
|
||||
else {
|
||||
else if (m_jobSize <= 256) {
|
||||
m_blake2b_initial_hash_double->enqueue(m_queue, m_intensity);
|
||||
}
|
||||
else {
|
||||
m_blake2b_initial_hash_big->enqueue(m_queue, m_intensity);
|
||||
}
|
||||
|
||||
m_fillAes1Rx4_scratchpad->enqueue(m_queue, m_intensity);
|
||||
|
||||
@ -150,12 +154,15 @@ void xmrig::OclRxBaseRunner::set(const Job &job, uint8_t *blob)
|
||||
memset(blob + job.size(), 0, Job::kMaxBlobSize - job.size());
|
||||
}
|
||||
|
||||
memset(blob + job.nonceOffset(), 0, job.nonceSize());
|
||||
|
||||
enqueueWriteBuffer(m_input, CL_TRUE, 0, Job::kMaxBlobSize, blob);
|
||||
|
||||
m_jobSize = job.size();
|
||||
|
||||
m_blake2b_initial_hash->setBlobSize(job.size());
|
||||
m_blake2b_initial_hash_double->setBlobSize(job.size());
|
||||
m_blake2b_initial_hash_big->setBlobSize(job.size());
|
||||
|
||||
m_find_shares->setTarget(job.target());
|
||||
}
|
||||
@ -191,6 +198,9 @@ void xmrig::OclRxBaseRunner::build()
|
||||
m_blake2b_initial_hash_double = new Blake2bInitialHashDoubleKernel(m_program);
|
||||
m_blake2b_initial_hash_double->setArgs(m_hashes, m_input);
|
||||
|
||||
m_blake2b_initial_hash_big = new Blake2bInitialHashBigKernel(m_program);
|
||||
m_blake2b_initial_hash_big->setArgs(m_hashes, m_input);
|
||||
|
||||
m_blake2b_hash_registers_32 = new Blake2bHashRegistersKernel(m_program, "blake2b_hash_registers_32");
|
||||
m_blake2b_hash_registers_64 = new Blake2bHashRegistersKernel(m_program, "blake2b_hash_registers_64");
|
||||
|
||||
|
||||
@ -36,6 +36,7 @@ namespace xmrig {
|
||||
class Blake2bHashRegistersKernel;
|
||||
class Blake2bInitialHashKernel;
|
||||
class Blake2bInitialHashDoubleKernel;
|
||||
class Blake2bInitialHashBigKernel;
|
||||
class FillAesKernel;
|
||||
class FindSharesKernel;
|
||||
class HashAesKernel;
|
||||
@ -53,7 +54,7 @@ protected:
|
||||
size_t bufferSize() const override;
|
||||
void build() override;
|
||||
void init() override;
|
||||
void run(uint32_t nonce, uint32_t *hashOutput) override;
|
||||
void run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput) override;
|
||||
void set(const Job &job, uint8_t *blob) override;
|
||||
|
||||
protected:
|
||||
@ -63,6 +64,7 @@ protected:
|
||||
Blake2bHashRegistersKernel *m_blake2b_hash_registers_64 = nullptr;
|
||||
Blake2bInitialHashKernel *m_blake2b_initial_hash = nullptr;
|
||||
Blake2bInitialHashDoubleKernel *m_blake2b_initial_hash_double = nullptr;
|
||||
Blake2bInitialHashBigKernel* m_blake2b_initial_hash_big = nullptr;
|
||||
Buffer m_seed;
|
||||
cl_mem m_dataset = nullptr;
|
||||
cl_mem m_entropy = nullptr;
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user