From e2d2591281b44a153c94c3686fed25984f0716ab Mon Sep 17 00:00:00 2001 From: XMRig Date: Sun, 1 Sep 2019 14:16:19 +0700 Subject: [PATCH] Implemented remaining kernels. --- src/backend/opencl/kernels/Cn2Kernel.cpp | 59 ++++++++++++++++++ src/backend/opencl/kernels/Cn2Kernel.h | 47 ++++++++++++++ src/backend/opencl/kernels/CnBranchKernel.cpp | 62 +++++++++++++++++++ src/backend/opencl/kernels/CnBranchKernel.h | 47 ++++++++++++++ src/backend/opencl/opencl.cmake | 4 ++ src/backend/opencl/runners/OclCnRunner.cpp | 46 ++++++++++++-- src/backend/opencl/runners/OclCnRunner.h | 15 +++-- 7 files changed, 270 insertions(+), 10 deletions(-) create mode 100644 src/backend/opencl/kernels/Cn2Kernel.cpp create mode 100644 src/backend/opencl/kernels/Cn2Kernel.h create mode 100644 src/backend/opencl/kernels/CnBranchKernel.cpp create mode 100644 src/backend/opencl/kernels/CnBranchKernel.h diff --git a/src/backend/opencl/kernels/Cn2Kernel.cpp b/src/backend/opencl/kernels/Cn2Kernel.cpp new file mode 100644 index 000000000..0e4a1d171 --- /dev/null +++ b/src/backend/opencl/kernels/Cn2Kernel.cpp @@ -0,0 +1,59 @@ +/* XMRig + * Copyright 2010 Jeff Garzik + * Copyright 2012-2014 pooler + * Copyright 2014 Lucas Jones + * Copyright 2014-2016 Wolf9466 + * Copyright 2016 Jay D Dee + * Copyright 2017-2018 XMR-Stak , + * Copyright 2018-2019 SChernykh + * Copyright 2016-2019 XMRig , + * + * 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 . + */ + + +#include "backend/opencl/kernels/Cn2Kernel.h" +#include "backend/opencl/wrappers/OclLib.h" + + +xmrig::Cn2Kernel::Cn2Kernel(cl_program program) : OclKernel(program, "cn2") +{ +} + + +bool xmrig::Cn2Kernel::enqueue(cl_command_queue queue, uint32_t nonce, size_t threads) +{ + const size_t offset[2] = { nonce, 1 }; + const size_t gthreads[2] = { threads, 8 }; + static const size_t lthreads[2] = { 8, 8 }; + + return enqueueNDRange(queue, 2, offset, gthreads, lthreads); +} + + +// __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global uint *Branch0, __global uint *Branch1, __global uint *Branch2, __global uint *Branch3, uint Threads) +bool xmrig::Cn2Kernel::setArgs(cl_mem scratchpads, cl_mem states, const std::vector &branches, uint32_t threads) +{ + if (!setArg(0, sizeof(cl_mem), &scratchpads) || !setArg(1, sizeof(cl_mem), &states) || !setArg(6, sizeof(uint32_t), &threads)) { + return false; + } + + for (uint32_t i = 0; i < branches.size(); ++i) { + if (!setArg(i + 2, sizeof(cl_mem), &branches[i])) { + return false; + } + } + + return true; +} diff --git a/src/backend/opencl/kernels/Cn2Kernel.h b/src/backend/opencl/kernels/Cn2Kernel.h new file mode 100644 index 000000000..a7b54ab2c --- /dev/null +++ b/src/backend/opencl/kernels/Cn2Kernel.h @@ -0,0 +1,47 @@ +/* XMRig + * Copyright 2010 Jeff Garzik + * Copyright 2012-2014 pooler + * Copyright 2014 Lucas Jones + * Copyright 2014-2016 Wolf9466 + * Copyright 2016 Jay D Dee + * Copyright 2017-2018 XMR-Stak , + * Copyright 2018-2019 SChernykh + * Copyright 2016-2019 XMRig , + * + * 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 . + */ + +#ifndef XMRIG_CN2KERNEL_H +#define XMRIG_CN2KERNEL_H + + +#include "backend/opencl/wrappers/OclKernel.h" + + +namespace xmrig { + + +class Cn2Kernel : public OclKernel +{ +public: + Cn2Kernel(cl_program program); + bool enqueue(cl_command_queue queue, uint32_t nonce, size_t threads); + bool setArgs(cl_mem scratchpads, cl_mem states, const std::vector &branches, uint32_t threads); +}; + + +} // namespace xmrig + + +#endif /* XMRIG_CN2KERNEL_H */ diff --git a/src/backend/opencl/kernels/CnBranchKernel.cpp b/src/backend/opencl/kernels/CnBranchKernel.cpp new file mode 100644 index 000000000..04dd830eb --- /dev/null +++ b/src/backend/opencl/kernels/CnBranchKernel.cpp @@ -0,0 +1,62 @@ +/* XMRig + * Copyright 2010 Jeff Garzik + * Copyright 2012-2014 pooler + * Copyright 2014 Lucas Jones + * Copyright 2014-2016 Wolf9466 + * Copyright 2016 Jay D Dee + * Copyright 2017-2018 XMR-Stak , + * Copyright 2018-2019 SChernykh + * Copyright 2016-2019 XMRig , + * + * 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 . + */ + + +#include "backend/opencl/kernels/CnBranchKernel.h" +#include "backend/opencl/wrappers/OclLib.h" + + +namespace xmrig { + + +static const char *names[4] = { "Blake", "Groestl", "JH", "Skein" }; + + +} // namespace xmrig + + +xmrig::CnBranchKernel::CnBranchKernel(size_t index, cl_program program) : OclKernel(program, names[index]) +{ +} + + +bool xmrig::CnBranchKernel::enqueue(cl_command_queue queue, uint32_t nonce, size_t threads, size_t worksize) +{ + const size_t offset = nonce; + const size_t gthreads = threads; + const size_t lthreads = worksize; + + return enqueueNDRange(queue, 1, &offset, >hreads, <hreads); +} + + +// __kernel void Skein(__global ulong *states, __global uint *BranchBuf, __global uint *output, ulong Target, uint Threads) +bool xmrig::CnBranchKernel::setArgs(cl_mem states, cl_mem branch, cl_mem output, uint64_t target, uint32_t threads) +{ + return setArg(0, sizeof(cl_mem), &states) && + setArg(1, sizeof(cl_mem), &branch) && + setArg(2, sizeof(cl_mem), &output) && + setArg(3, sizeof(cl_ulong), &target) && + setArg(4, sizeof(cl_uint), &threads); +} diff --git a/src/backend/opencl/kernels/CnBranchKernel.h b/src/backend/opencl/kernels/CnBranchKernel.h new file mode 100644 index 000000000..496d7d2d0 --- /dev/null +++ b/src/backend/opencl/kernels/CnBranchKernel.h @@ -0,0 +1,47 @@ +/* XMRig + * Copyright 2010 Jeff Garzik + * Copyright 2012-2014 pooler + * Copyright 2014 Lucas Jones + * Copyright 2014-2016 Wolf9466 + * Copyright 2016 Jay D Dee + * Copyright 2017-2018 XMR-Stak , + * Copyright 2018-2019 SChernykh + * Copyright 2016-2019 XMRig , + * + * 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 . + */ + +#ifndef XMRIG_CNBRANCHKERNEL_H +#define XMRIG_CNBRANCHKERNEL_H + + +#include "backend/opencl/wrappers/OclKernel.h" + + +namespace xmrig { + + +class CnBranchKernel : public OclKernel +{ +public: + CnBranchKernel(size_t index, cl_program program); + bool enqueue(cl_command_queue queue, uint32_t nonce, size_t threads, size_t worksize); + bool setArgs(cl_mem states, cl_mem branch, cl_mem output, uint64_t target, uint32_t threads); +}; + + +} // namespace xmrig + + +#endif /* XMRIG_CNBRANCHKERNEL_H */ diff --git a/src/backend/opencl/opencl.cmake b/src/backend/opencl/opencl.cmake index da89ebf4e..6c94cdc1c 100644 --- a/src/backend/opencl/opencl.cmake +++ b/src/backend/opencl/opencl.cmake @@ -8,6 +8,8 @@ if (WITH_OPENCL) src/backend/opencl/interfaces/IOclRunner.h src/backend/opencl/kernels/Cn0Kernel.h src/backend/opencl/kernels/Cn1Kernel.h + src/backend/opencl/kernels/Cn2Kernel.h + src/backend/opencl/kernels/CnBranchKernel.h src/backend/opencl/OclBackend.h src/backend/opencl/OclCache.h src/backend/opencl/OclConfig.h @@ -30,6 +32,8 @@ if (WITH_OPENCL) src/backend/opencl/cl/OclSource.cpp src/backend/opencl/kernels/Cn0Kernel.cpp src/backend/opencl/kernels/Cn1Kernel.cpp + src/backend/opencl/kernels/Cn2Kernel.cpp + src/backend/opencl/kernels/CnBranchKernel.cpp src/backend/opencl/OclBackend.cpp src/backend/opencl/OclCache.cpp src/backend/opencl/OclConfig.cpp diff --git a/src/backend/opencl/runners/OclCnRunner.cpp b/src/backend/opencl/runners/OclCnRunner.cpp index da06a5886..88d2c72f7 100644 --- a/src/backend/opencl/runners/OclCnRunner.cpp +++ b/src/backend/opencl/runners/OclCnRunner.cpp @@ -25,6 +25,8 @@ #include "backend/opencl/kernels/Cn0Kernel.h" #include "backend/opencl/kernels/Cn1Kernel.h" +#include "backend/opencl/kernels/Cn2Kernel.h" +#include "backend/opencl/kernels/CnBranchKernel.h" #include "backend/opencl/OclLaunchData.h" #include "backend/opencl/runners/OclCnRunner.h" #include "backend/opencl/wrappers/OclLib.h" @@ -89,11 +91,14 @@ xmrig::OclCnRunner::OclCnRunner(size_t index, const OclLaunchData &data) : OclBa xmrig::OclCnRunner::~OclCnRunner() { delete m_cn0; + delete m_cn1; + delete m_cn2; OclLib::releaseMemObject(m_scratchpads); OclLib::releaseMemObject(m_states); for (size_t i = 0; i < BRANCH_MAX; ++i) { + delete m_branchKernels[i]; OclLib::releaseMemObject(m_branches[i]); } } @@ -115,9 +120,6 @@ bool xmrig::OclCnRunner::run(uint32_t nonce, uint32_t *hashOutput) { static const cl_uint zero = 0; - cl_int ret; - size_t branchNonces[4] = { 0 }; - const size_t g_intensity = data().thread.intensity(); const size_t w_size = data().thread.worksize(); const size_t g_thd = ((g_intensity + w_size - 1u) / w_size) * w_size; @@ -142,7 +144,24 @@ bool xmrig::OclCnRunner::run(uint32_t nonce, uint32_t *hashOutput) return false; } - OclLib::finish(m_queue); + if (!m_cn2->enqueue(m_queue, nonce, g_thd)) { + return false; + } + + for (size_t i = 0; i < BRANCH_MAX; ++i) { + if (!m_branchKernels[i]->enqueue(m_queue, nonce, g_thd, w_size)) { + return false; + } + } + + if (OclLib::enqueueReadBuffer(m_queue, m_output, CL_TRUE, 0, sizeof(cl_uint) * 0x100, hashOutput, 0, nullptr, nullptr) != CL_SUCCESS) { + return false; + } + + uint32_t &results = hashOutput[0xFF]; + if (results > 0xFF) { + results = 0xFF; + } return true; } @@ -152,7 +171,7 @@ bool xmrig::OclCnRunner::selfTest() const { return OclBaseRunner::selfTest() && m_cn0->isValid() && - m_cn1->isValid(); + m_cn1->isValid() && m_cn2->isValid(); } @@ -177,6 +196,18 @@ bool xmrig::OclCnRunner::set(const Job &job, uint8_t *blob) return false; } + const uint32_t intensity = data().thread.intensity(); + + if (!m_cn2->setArgs(m_scratchpads, m_states, m_branches, intensity)) { + return false; + } + + for (size_t i = 0; i < BRANCH_MAX; ++i) { + if (!m_branchKernels[i]->setArgs(m_states, m_branches[i], m_output, job.target(), intensity)) { + return false; + } + } + return true; } @@ -191,4 +222,9 @@ void xmrig::OclCnRunner::build() m_cn0 = new Cn0Kernel(m_program); m_cn1 = new Cn1Kernel(m_program); + m_cn2 = new Cn2Kernel(m_program); + + for (size_t i = 0; i < BRANCH_MAX; ++i) { + m_branchKernels[i] = new CnBranchKernel(i, m_program); + } } diff --git a/src/backend/opencl/runners/OclCnRunner.h b/src/backend/opencl/runners/OclCnRunner.h index 8ad4c7e5b..2afa44eaf 100644 --- a/src/backend/opencl/runners/OclCnRunner.h +++ b/src/backend/opencl/runners/OclCnRunner.h @@ -34,6 +34,8 @@ namespace xmrig { class Cn0Kernel; class Cn1Kernel; +class Cn2Kernel; +class CnBranchKernel; class OclCnRunner : public OclBaseRunner @@ -59,11 +61,14 @@ private: }; - cl_mem m_branches[BRANCH_MAX] = { nullptr, nullptr, nullptr, nullptr }; - cl_mem m_scratchpads = nullptr; - cl_mem m_states = nullptr; - Cn0Kernel *m_cn0 = nullptr; - Cn1Kernel *m_cn1 = nullptr; + cl_mem m_scratchpads = nullptr; + cl_mem m_states = nullptr; + Cn0Kernel *m_cn0 = nullptr; + Cn1Kernel *m_cn1 = nullptr; + Cn2Kernel *m_cn2 = nullptr; + + std::vector m_branches = { nullptr, nullptr, nullptr, nullptr }; + std::vector m_branchKernels = { nullptr, nullptr, nullptr, nullptr }; };