diff --git a/src/backend/opencl/OclWorker.cpp b/src/backend/opencl/OclWorker.cpp index 4bc2654ab..06f999cf8 100644 --- a/src/backend/opencl/OclWorker.cpp +++ b/src/backend/opencl/OclWorker.cpp @@ -91,6 +91,8 @@ bool xmrig::OclWorker::selfTest() void xmrig::OclWorker::start() { + cl_uint results[0x100]; + while (Nonce::sequence(Nonce::OPENCL) > 0) { if (Nonce::isPaused()) { do { @@ -116,6 +118,7 @@ void xmrig::OclWorker::start() break; } + m_runner->run(results); std::this_thread::sleep_for(std::chrono::milliseconds(2000)); // FIXME m_job.nextRound(kReserveCount); diff --git a/src/backend/opencl/kernels/Cn0Kernel.cpp b/src/backend/opencl/kernels/Cn0Kernel.cpp new file mode 100644 index 000000000..a7674bc76 --- /dev/null +++ b/src/backend/opencl/kernels/Cn0Kernel.cpp @@ -0,0 +1,42 @@ +/* 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/Cn0Kernel.h" +#include "backend/opencl/wrappers/OclLib.h" + + +xmrig::Cn0Kernel::Cn0Kernel(cl_program program) : OclKernel(program, "cn0") +{ +} + + +// __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states, uint Threads) +bool xmrig::Cn0Kernel::setArgs(cl_mem input, cl_mem scratchpads, cl_mem states, uint32_t threads) +{ + return setArg(0, sizeof(cl_mem), &input) && + setArg(1, sizeof(cl_mem), &scratchpads) && + setArg(2, sizeof(cl_mem), &states) && + setArg(3, sizeof(uint32_t), &threads); +} diff --git a/src/backend/opencl/kernels/Cn0Kernel.h b/src/backend/opencl/kernels/Cn0Kernel.h new file mode 100644 index 000000000..5feda6acc --- /dev/null +++ b/src/backend/opencl/kernels/Cn0Kernel.h @@ -0,0 +1,49 @@ +/* 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_CN0KERNEL_H +#define XMRIG_CN0KERNEL_H + + +#include "backend/opencl/wrappers/OclKernel.h" + + +typedef struct _cl_mem *cl_mem; + + +namespace xmrig { + + +class Cn0Kernel : public OclKernel +{ +public: + Cn0Kernel(cl_program program); + bool setArgs(cl_mem input, cl_mem scratchpads, cl_mem states, uint32_t threads); +}; + + +} // namespace xmrig + + +#endif /* XMRIG_CN0KERNEL_H */ diff --git a/src/backend/opencl/opencl.cmake b/src/backend/opencl/opencl.cmake index 543354d48..befccde8e 100644 --- a/src/backend/opencl/opencl.cmake +++ b/src/backend/opencl/opencl.cmake @@ -6,6 +6,7 @@ if (WITH_OPENCL) set(HEADERS_BACKEND_OPENCL src/backend/opencl/cl/OclSource.h src/backend/opencl/interfaces/IOclRunner.h + src/backend/opencl/kernels/Cn0Kernel.h src/backend/opencl/OclBackend.h src/backend/opencl/OclCache.h src/backend/opencl/OclConfig.h @@ -18,6 +19,7 @@ if (WITH_OPENCL) src/backend/opencl/wrappers/OclContext.h src/backend/opencl/wrappers/OclDevice.h src/backend/opencl/wrappers/OclError.h + src/backend/opencl/wrappers/OclKernel.h src/backend/opencl/wrappers/OclLib.h src/backend/opencl/wrappers/OclPlatform.h src/backend/opencl/wrappers/OclVendor.h @@ -25,6 +27,7 @@ if (WITH_OPENCL) set(SOURCES_BACKEND_OPENCL src/backend/opencl/cl/OclSource.cpp + src/backend/opencl/kernels/Cn0Kernel.cpp src/backend/opencl/OclBackend.cpp src/backend/opencl/OclCache.cpp src/backend/opencl/OclConfig.cpp @@ -37,6 +40,7 @@ if (WITH_OPENCL) src/backend/opencl/wrappers/OclContext.cpp src/backend/opencl/wrappers/OclDevice.cpp src/backend/opencl/wrappers/OclError.cpp + src/backend/opencl/wrappers/OclKernel.cpp src/backend/opencl/wrappers/OclLib.cpp src/backend/opencl/wrappers/OclPlatform.cpp ) diff --git a/src/backend/opencl/runners/OclCnRunner.cpp b/src/backend/opencl/runners/OclCnRunner.cpp index e62cdd17a..e72132c68 100644 --- a/src/backend/opencl/runners/OclCnRunner.cpp +++ b/src/backend/opencl/runners/OclCnRunner.cpp @@ -23,6 +23,7 @@ */ +#include "backend/opencl/kernels/Cn0Kernel.h" #include "backend/opencl/OclLaunchData.h" #include "backend/opencl/runners/OclCnRunner.h" #include "backend/opencl/wrappers/OclLib.h" @@ -80,6 +81,8 @@ xmrig::OclCnRunner::OclCnRunner(size_t index, const OclLaunchData &data) : OclBa xmrig::OclCnRunner::~OclCnRunner() { + delete m_cn0; + OclLib::releaseMemObject(m_scratchpads); OclLib::releaseMemObject(m_states); OclLib::releaseMemObject(m_blake256); @@ -101,6 +104,12 @@ bool xmrig::OclCnRunner::isReadyToBuild() const } +bool xmrig::OclCnRunner::selfTest() const +{ + return OclBaseRunner::selfTest() && m_cn0->isValid(); +} + + bool xmrig::OclCnRunner::set(const Job &job, uint8_t *blob) { if (job.size() > (Job::kMaxBlobSize - 4)) { @@ -114,5 +123,22 @@ bool xmrig::OclCnRunner::set(const Job &job, uint8_t *blob) return false; } + if (!m_cn0->setArgs(m_input, m_scratchpads, m_states, data().thread.intensity())) { + return false; + } + + LOG_WARN(GREEN_S "OK"); return false; } + + +void xmrig::OclCnRunner::build() +{ + OclBaseRunner::build(); + + if (!m_program) { + return; + } + + m_cn0 = new Cn0Kernel(m_program); +} diff --git a/src/backend/opencl/runners/OclCnRunner.h b/src/backend/opencl/runners/OclCnRunner.h index ac3691e3c..c23e1c1e7 100644 --- a/src/backend/opencl/runners/OclCnRunner.h +++ b/src/backend/opencl/runners/OclCnRunner.h @@ -32,6 +32,9 @@ namespace xmrig { +class Cn0Kernel; + + class OclCnRunner : public OclBaseRunner { public: @@ -40,7 +43,9 @@ public: protected: bool isReadyToBuild() const override; + bool selfTest() const override; bool set(const Job &job, uint8_t *blob) override; + void build() override; private: cl_mem m_blake256 = nullptr; @@ -49,6 +54,7 @@ private: cl_mem m_scratchpads = nullptr; cl_mem m_skein512 = nullptr; cl_mem m_states = nullptr; + Cn0Kernel *m_cn0 = nullptr; }; diff --git a/src/backend/opencl/wrappers/OclKernel.cpp b/src/backend/opencl/wrappers/OclKernel.cpp new file mode 100644 index 000000000..6df687a2b --- /dev/null +++ b/src/backend/opencl/wrappers/OclKernel.cpp @@ -0,0 +1,60 @@ +/* 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/wrappers/OclError.h" +#include "backend/opencl/wrappers/OclKernel.h" +#include "backend/opencl/wrappers/OclLib.h" +#include "base/io/log/Log.h" + + +xmrig::OclKernel::OclKernel(cl_program program, const char *name) : + m_name(name) +{ + cl_int ret = 0; + m_kernel = OclLib::createKernel(program, name, &ret); +} + + +xmrig::OclKernel::~OclKernel() +{ + OclLib::releaseKernel(m_kernel); +} + + +bool xmrig::OclKernel::setArg(uint32_t index, size_t size, const void *value) +{ + if (!isValid()) { + return false; + } + + const cl_int ret = OclLib::setKernelArg(m_kernel, index, size, value); + if (ret != CL_SUCCESS) { + LOG_ERR(MAGENTA_BG_BOLD(WHITE_BOLD_S " ocl ") RED(" error ") RED_BOLD("%s") RED(" when calling ") RED_BOLD("clSetKernelArg") RED(" for kernel ") RED_BOLD("%s") + RED(" argument ") RED_BOLD("%u") RED(" size ") RED_BOLD("%zu"), + OclError::toString(ret), name().data(), index, size); + } + + return ret == CL_SUCCESS; +} diff --git a/src/backend/opencl/wrappers/OclKernel.h b/src/backend/opencl/wrappers/OclKernel.h new file mode 100644 index 000000000..bc7f8d071 --- /dev/null +++ b/src/backend/opencl/wrappers/OclKernel.h @@ -0,0 +1,60 @@ +/* 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_OCLKERNEL_H +#define XMRIG_OCLKERNEL_H + + +#include "base/tools/String.h" + + +typedef struct _cl_kernel *cl_kernel; +typedef struct _cl_program *cl_program; + + +namespace xmrig { + + +class OclKernel +{ +public: + OclKernel(cl_program program, const char *name); + virtual ~OclKernel(); + + inline bool isValid() const { return m_kernel != nullptr; } + inline cl_kernel kernel() const { return m_kernel; } + inline const String &name() const { return m_name; } + + bool setArg(uint32_t index, size_t size, const void *value); + +private: + cl_kernel m_kernel = nullptr; + const String m_name; +}; + + +} // namespace xmrig + + +#endif /* XMRIG_OCLKERNEL_H */ diff --git a/src/backend/opencl/wrappers/OclLib.cpp b/src/backend/opencl/wrappers/OclLib.cpp index d536db74a..e2f6b0651 100644 --- a/src/backend/opencl/wrappers/OclLib.cpp +++ b/src/backend/opencl/wrappers/OclLib.cpp @@ -469,7 +469,10 @@ cl_kernel xmrig::OclLib::createKernel(cl_program program, const char *kernel_nam auto result = pCreateKernel(program, kernel_name, errcode_ret); if (*errcode_ret != CL_SUCCESS) { - LOG_ERR("Error %s when calling clCreateKernel for kernel %s.", OclError::toString(*errcode_ret), kernel_name); + LOG_ERR(MAGENTA_BG_BOLD(WHITE_BOLD_S " ocl ") RED(" error ") RED_BOLD("%s") RED(" when calling ") RED_BOLD("clCreateKernel") RED(" for kernel ") RED_BOLD("%s"), + OclError::toString(*errcode_ret), kernel_name); + + return nullptr; } return result;