diff --git a/src/backend/opencl/OclThread.cpp b/src/backend/opencl/OclThread.cpp index 70fd5725..24c4ffb4 100644 --- a/src/backend/opencl/OclThread.cpp +++ b/src/backend/opencl/OclThread.cpp @@ -73,7 +73,7 @@ xmrig::OclThread::OclThread(const rapidjson::Value &value) m_stridedIndex = stridedIndex.GetBool() ? 1 : 0; } else if (stridedIndex.IsUint()) { - m_stridedIndex = std::max(stridedIndex.GetUint(), 2u); + m_stridedIndex = std::min(stridedIndex.GetUint(), 2u); } } @@ -131,5 +131,5 @@ rapidjson::Value xmrig::OclThread::toJSON(rapidjson::Document &doc) const void xmrig::OclThread::setUnrollFactor(uint32_t unrollFactor) { - m_unrollFactor = unrollFactor == 0 ? 1 : std::max(unrollFactor, 128u); + m_unrollFactor = unrollFactor == 0 ? 1 : std::min(unrollFactor, 128u); } diff --git a/src/backend/opencl/OclThread.h b/src/backend/opencl/OclThread.h index d61b5364..3c6a98c9 100644 --- a/src/backend/opencl/OclThread.h +++ b/src/backend/opencl/OclThread.h @@ -50,7 +50,6 @@ public: inline bool isCompMode() const { return m_compMode; } inline bool isValid() const { return m_intensity > 0; } inline int64_t affinity() const { return m_affinity; } - inline int64_t worksize() const { return m_worksize; } inline uint32_t bfactor() const { return m_bfactor; } inline uint32_t datasetHost() const { return m_datasetHost < 0 ? 0 : static_cast(m_datasetHost); } inline uint32_t gcnAsm() const { return m_gcnAsm; } @@ -59,6 +58,7 @@ public: inline uint32_t memChunk() const { return m_memChunk; } inline uint32_t stridedIndex() const { return m_stridedIndex; } inline uint32_t unrollFactor() const { return m_unrollFactor; } + inline uint32_t worksize() const { return m_worksize; } inline bool operator!=(const OclThread &other) const { return !isEqual(other); } inline bool operator==(const OclThread &other) const { return isEqual(other); } diff --git a/src/backend/opencl/OclWorker.cpp b/src/backend/opencl/OclWorker.cpp index 4d349f38..f9e9d900 100644 --- a/src/backend/opencl/OclWorker.cpp +++ b/src/backend/opencl/OclWorker.cpp @@ -29,11 +29,17 @@ #include "backend/opencl/OclWorker.h" +#include "backend/opencl/runners/OclCnRunner.h" #include "core/Miner.h" #include "crypto/common/Nonce.h" #include "net/JobResults.h" +#ifdef XMRIG_ALGO_RANDOMX +# include "backend/opencl/runners/OclRxRunner.h" +#endif + + namespace xmrig { static constexpr uint32_t kReserveCount = 4096; @@ -47,17 +53,35 @@ xmrig::OclWorker::OclWorker(size_t index, const OclLaunchData &data) : m_algorithm(data.algorithm), m_miner(data.miner) { + switch (m_algorithm.family()) { + case Algorithm::RANDOM_X: +# ifdef XMRIG_ALGO_RANDOMX + m_runner = new OclRxRunner(index, data); +# endif + break; + + case Algorithm::ARGON2: +# ifdef XMRIG_ALGO_ARGON2 + m_runner = nullptr; // TODO +# endif + break; + + default: + m_runner = new OclCnRunner(index, data); + break; + } } xmrig::OclWorker::~OclWorker() { + delete m_runner; } bool xmrig::OclWorker::selfTest() { - return true; + return m_runner && m_runner->selfTest(); } diff --git a/src/backend/opencl/OclWorker.h b/src/backend/opencl/OclWorker.h index cbe50524..0f73e13b 100644 --- a/src/backend/opencl/OclWorker.h +++ b/src/backend/opencl/OclWorker.h @@ -36,6 +36,9 @@ namespace xmrig { +class IOclRunner; + + class OclWorker : public Worker { public: @@ -51,6 +54,7 @@ private: const Algorithm m_algorithm; const Miner *m_miner; + IOclRunner *m_runner = nullptr; WorkerJob<1> m_job; }; diff --git a/src/backend/opencl/interfaces/IOclRunner.h b/src/backend/opencl/interfaces/IOclRunner.h new file mode 100644 index 00000000..efe349f2 --- /dev/null +++ b/src/backend/opencl/interfaces/IOclRunner.h @@ -0,0 +1,53 @@ +/* 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_IOCLRUNNER_H +#define XMRIG_IOCLRUNNER_H + + +#include + + +namespace xmrig { + + +class Job; + + +class IOclRunner +{ +public: + virtual ~IOclRunner() = default; + + virtual bool selfTest() const = 0; + virtual const char *buildOptions() const = 0; + virtual void run(uint32_t *hashOutput) = 0; + virtual void set(const Job &job) = 0; +}; + + +} /* namespace xmrig */ + + +#endif // XMRIG_IOCLRUNNER_H diff --git a/src/backend/opencl/opencl.cmake b/src/backend/opencl/opencl.cmake index 8b68d2c3..bea54f51 100644 --- a/src/backend/opencl/opencl.cmake +++ b/src/backend/opencl/opencl.cmake @@ -5,6 +5,7 @@ if (WITH_OPENCL) set(HEADERS_BACKEND_OPENCL src/backend/opencl/cl/OclSource.h + src/backend/opencl/interfaces/IOclRunner.h src/backend/opencl/OclBackend.h src/backend/opencl/OclCache.h src/backend/opencl/OclConfig.h @@ -12,6 +13,8 @@ if (WITH_OPENCL) src/backend/opencl/OclThread.h src/backend/opencl/OclThreads.h src/backend/opencl/OclWorker.h + src/backend/opencl/runners/OclBaseRunner.h + src/backend/opencl/runners/OclCnRunner.h src/backend/opencl/wrappers/OclContext.h src/backend/opencl/wrappers/OclDevice.h src/backend/opencl/wrappers/OclError.h @@ -29,6 +32,8 @@ if (WITH_OPENCL) src/backend/opencl/OclThread.cpp src/backend/opencl/OclThreads.cpp src/backend/opencl/OclWorker.cpp + src/backend/opencl/runners/OclBaseRunner.cpp + src/backend/opencl/runners/OclCnRunner.cpp src/backend/opencl/wrappers/OclContext.cpp src/backend/opencl/wrappers/OclDevice.cpp src/backend/opencl/wrappers/OclError.cpp @@ -41,6 +46,11 @@ if (WITH_OPENCL) else() list(APPEND SOURCES_BACKEND_OPENCL src/backend/opencl/OclCache_unix.cpp) endif() + + if (WITH_RANDOMX) + list(APPEND HEADERS_BACKEND_OPENCL src/backend/opencl/runners/OclRxRunner.h) + list(APPEND SOURCES_BACKEND_OPENCL src/backend/opencl/runners/OclRxRunner.cpp) + endif() else() remove_definitions(/DXMRIG_FEATURE_OPENCL) diff --git a/src/backend/opencl/runners/OclBaseRunner.cpp b/src/backend/opencl/runners/OclBaseRunner.cpp new file mode 100644 index 00000000..ee9a8fe8 --- /dev/null +++ b/src/backend/opencl/runners/OclBaseRunner.cpp @@ -0,0 +1,78 @@ +/* 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/OclLaunchData.h" +#include "backend/opencl/runners/OclBaseRunner.h" +#include "backend/opencl/wrappers/OclLib.h" +#include "base/net/stratum/Job.h" + + +xmrig::OclBaseRunner::OclBaseRunner(size_t, const OclLaunchData &data) : + m_algorithm(data.algorithm), + m_ctx(data.ctx) +{ + cl_int ret; + m_queue = OclLib::createCommandQueue(m_ctx, data.device.id(), &ret); + if (ret != CL_SUCCESS) { + return; + } + + m_input = OclLib::createBuffer(m_ctx, CL_MEM_READ_ONLY, Job::kMaxBlobSize, nullptr, &ret); + m_output = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, sizeof(cl_uint) * 0x100, nullptr, &ret); +} + + +xmrig::OclBaseRunner::~OclBaseRunner() +{ + OclLib::releaseMemObject(m_input); + OclLib::releaseMemObject(m_output); + + OclLib::releaseCommandQueue(m_queue); +} + + +bool xmrig::OclBaseRunner::selfTest() const +{ + return m_queue != nullptr && m_input != nullptr && m_output != nullptr && !m_options.empty(); +} + + + +const char *xmrig::OclBaseRunner::buildOptions() const +{ + return m_options.c_str(); +} + + +void xmrig::OclBaseRunner::run(uint32_t *hashOutput) +{ + +} + + +void xmrig::OclBaseRunner::set(const Job &job) +{ + +} diff --git a/src/backend/opencl/runners/OclBaseRunner.h b/src/backend/opencl/runners/OclBaseRunner.h new file mode 100644 index 00000000..2b8fa791 --- /dev/null +++ b/src/backend/opencl/runners/OclBaseRunner.h @@ -0,0 +1,68 @@ +/* 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_OCLBASERUNNER_H +#define XMRIG_OCLBASERUNNER_H + + +#include + + +#include "3rdparty/cl.h" +#include "backend/opencl/interfaces/IOclRunner.h" +#include "crypto/common/Algorithm.h" + + +namespace xmrig { + + +class OclLaunchData; + + +class OclBaseRunner : public IOclRunner +{ +public: + OclBaseRunner(size_t index, const OclLaunchData &data); + ~OclBaseRunner() override; + +protected: + bool selfTest() const override; + const char *buildOptions() const override; + void run(uint32_t *hashOutput) override; + void set(const Job &job) override; + +protected: + Algorithm m_algorithm; + cl_command_queue m_queue = nullptr; + cl_context m_ctx; + cl_mem m_input = nullptr; + cl_mem m_output = nullptr; + std::string m_options; +}; + + +} /* namespace xmrig */ + + +#endif // XMRIG_OCLBASERUNNER_H diff --git a/src/backend/opencl/runners/OclCnRunner.cpp b/src/backend/opencl/runners/OclCnRunner.cpp new file mode 100644 index 00000000..8c87309e --- /dev/null +++ b/src/backend/opencl/runners/OclCnRunner.cpp @@ -0,0 +1,99 @@ +/* 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/runners/OclCnRunner.h" +#include "backend/opencl/wrappers/OclLib.h" +#include "backend/opencl/OclLaunchData.h" +#include "crypto/cn/CnAlgo.h" + + +xmrig::OclCnRunner::OclCnRunner(size_t index, const OclLaunchData &data) : OclBaseRunner(index, data) +{ + if (m_queue == nullptr) { + return; + } + + const size_t g_thd = data.thread.intensity(); + + cl_int ret; + m_scratchpads = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, data.algorithm.l3() * g_thd, nullptr, &ret); + if (ret != CL_SUCCESS) { + return; + } + + m_states = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, 200 * g_thd, nullptr, &ret); + m_blake256 = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, sizeof(cl_uint) * (g_thd + 2), nullptr, &ret); + m_groestl256 = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, sizeof(cl_uint) * (g_thd + 2), nullptr, &ret); + m_jh256 = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, sizeof(cl_uint) * (g_thd + 2), nullptr, &ret); + m_skein512 = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, sizeof(cl_uint) * (g_thd + 2), nullptr, &ret); + + uint32_t stridedIndex = data.thread.stridedIndex(); + if (data.device.vendorId() == OCL_VENDOR_NVIDIA) { + stridedIndex = 0; + } + else if (stridedIndex == 1 && (m_algorithm.family() == Algorithm::CN_PICO || (m_algorithm.family() == Algorithm::CN && CnAlgo<>::base(m_algorithm) == Algorithm::CN_2))) { + stridedIndex = 2; + } + + m_options += " -DITERATIONS=" + std::to_string(CnAlgo<>::iterations(m_algorithm)) + "U"; + m_options += " -DMASK=" + std::to_string(CnAlgo<>::mask(m_algorithm)) + "U"; + m_options += " -DWORKSIZE=" + std::to_string(data.thread.worksize()) + "U"; + m_options += " -DSTRIDED_INDEX=" + std::to_string(stridedIndex) + "U"; + m_options += " -DMEM_CHUNK_EXPONENT=" + std::to_string(1u << data.thread.memChunk()) + "U"; + m_options += " -DCOMP_MODE=" + std::to_string(data.thread.isCompMode() && g_thd % data.thread.worksize() != 0 ? 1u : 0u) + "U"; + m_options += " -DMEMORY=" + std::to_string(m_algorithm.l3()) + "LU"; + m_options += " -DALGO=" + std::to_string(m_algorithm.id()); + m_options += " -DALGO_FAMILY=" + std::to_string(m_algorithm.family()); + m_options += " -DCN_UNROLL=" + std::to_string(data.thread.unrollFactor()); + +# ifdef XMRIG_ALGO_CN_GPU + if (data.algorithm == Algorithm::CN_GPU) { + m_options += " -cl-fp32-correctly-rounded-divide-sqrt"; + } +# endif +} + + +xmrig::OclCnRunner::~OclCnRunner() +{ + OclLib::releaseMemObject(m_scratchpads); + OclLib::releaseMemObject(m_states); + OclLib::releaseMemObject(m_blake256); + OclLib::releaseMemObject(m_groestl256); + OclLib::releaseMemObject(m_jh256); + OclLib::releaseMemObject(m_skein512); +} + + +bool xmrig::OclCnRunner::selfTest() const +{ + return OclBaseRunner::selfTest() && + m_scratchpads != nullptr && + m_states != nullptr && + m_blake256 != nullptr && + m_groestl256 != nullptr && + m_jh256 != nullptr && + m_skein512 != nullptr; +} diff --git a/src/backend/opencl/runners/OclCnRunner.h b/src/backend/opencl/runners/OclCnRunner.h new file mode 100644 index 00000000..31f22c89 --- /dev/null +++ b/src/backend/opencl/runners/OclCnRunner.h @@ -0,0 +1,57 @@ +/* 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_OCLCNRUNNER_H +#define XMRIG_OCLCNRUNNER_H + + +#include "backend/opencl/runners/OclBaseRunner.h" + + +namespace xmrig { + + +class OclCnRunner : public OclBaseRunner +{ +public: + OclCnRunner(size_t index, const OclLaunchData &data); + ~OclCnRunner() override; + +protected: + bool selfTest() const override; + +private: + cl_mem m_blake256 = nullptr; + cl_mem m_groestl256 = nullptr; + cl_mem m_jh256 = nullptr; + cl_mem m_scratchpads = nullptr; + cl_mem m_skein512 = nullptr; + cl_mem m_states = nullptr; +}; + + +} /* namespace xmrig */ + + +#endif // XMRIG_OCLCNRUNNER_H diff --git a/src/backend/opencl/runners/OclRxRunner.cpp b/src/backend/opencl/runners/OclRxRunner.cpp new file mode 100644 index 00000000..6378defe --- /dev/null +++ b/src/backend/opencl/runners/OclRxRunner.cpp @@ -0,0 +1,36 @@ +/* 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/runners/OclRxRunner.h" + + +xmrig::OclRxRunner::OclRxRunner(size_t index, const OclLaunchData &data) : OclBaseRunner(index, data) +{ +} + + +bool xmrig::OclRxRunner::selfTest() const +{ + return false; // TODO +} diff --git a/src/backend/opencl/runners/OclRxRunner.h b/src/backend/opencl/runners/OclRxRunner.h new file mode 100644 index 00000000..09693555 --- /dev/null +++ b/src/backend/opencl/runners/OclRxRunner.h @@ -0,0 +1,48 @@ +/* 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_OCLRXRUNNER_H +#define XMRIG_OCLRXRUNNER_H + + +#include "backend/opencl/runners/OclBaseRunner.h" + + +namespace xmrig { + + +class OclRxRunner : public OclBaseRunner +{ +public: + OclRxRunner(size_t index, const OclLaunchData &data); + +protected: + bool selfTest() const override; +}; + + +} /* namespace xmrig */ + + +#endif // XMRIG_OCLRXRUNNER_H diff --git a/src/backend/opencl/wrappers/OclLib.cpp b/src/backend/opencl/wrappers/OclLib.cpp index 1fe56e07..3aa5fc25 100644 --- a/src/backend/opencl/wrappers/OclLib.cpp +++ b/src/backend/opencl/wrappers/OclLib.cpp @@ -23,6 +23,7 @@ */ +#include #include @@ -47,6 +48,7 @@ static const char *kEnqueueNDRangeKernel = "clEnqueueNDRangeKernel"; static const char *kEnqueueReadBuffer = "clEnqueueReadBuffer"; static const char *kEnqueueWriteBuffer = "clEnqueueWriteBuffer"; static const char *kFinish = "clFinish"; +static const char *kGetCommandQueueInfo = "clGetCommandQueueInfo"; static const char *kGetDeviceIDs = "clGetDeviceIDs"; static const char *kGetDeviceInfo = "clGetDeviceInfo"; static const char *kGetKernelInfo = "clGetKernelInfo"; @@ -72,6 +74,7 @@ typedef cl_int (CL_API_CALL *enqueueNDRangeKernel_t)(cl_command_queue, cl_kernel typedef cl_int (CL_API_CALL *enqueueReadBuffer_t)(cl_command_queue, cl_mem, cl_bool, size_t, size_t, void *, cl_uint, const cl_event *, cl_event *); typedef cl_int (CL_API_CALL *enqueueWriteBuffer_t)(cl_command_queue, cl_mem, cl_bool, size_t, size_t, const void *, cl_uint, const cl_event *, cl_event *); typedef cl_int (CL_API_CALL *finish_t)(cl_command_queue); +typedef cl_int (CL_API_CALL *getCommandQueueInfo_t)(cl_command_queue, cl_command_queue_info, size_t, void *, size_t *); typedef cl_int (CL_API_CALL *getDeviceIDs_t)(cl_platform_id, cl_device_type, cl_uint, cl_device_id *, cl_uint *); typedef cl_int (CL_API_CALL *getDeviceInfo_t)(cl_device_id, cl_device_info, size_t, void *, size_t *); typedef cl_int (CL_API_CALL *getKernelInfo_t)(cl_kernel, cl_kernel_info, size_t, void *, size_t *); @@ -106,6 +109,7 @@ static enqueueNDRangeKernel_t pEnqueueNDRangeKernel = nu static enqueueReadBuffer_t pEnqueueReadBuffer = nullptr; static enqueueWriteBuffer_t pEnqueueWriteBuffer = nullptr; static finish_t pFinish = nullptr; +static getCommandQueueInfo_t pGetCommandQueueInfo = nullptr; static getDeviceIDs_t pGetDeviceIDs = nullptr; static getDeviceInfo_t pGetDeviceInfo = nullptr; static getKernelInfo_t pGetKernelInfo = nullptr; @@ -182,6 +186,7 @@ bool xmrig::OclLib::load() DLSYM(ReleaseCommandQueue); DLSYM(ReleaseContext); DLSYM(GetKernelInfo); + DLSYM(GetCommandQueueInfo); # if defined(CL_VERSION_2_0) uv_dlsym(&oclLib, kCreateCommandQueueWithProperties, reinterpret_cast(&pCreateCommandQueueWithProperties)); @@ -222,6 +227,8 @@ cl_command_queue xmrig::OclLib::createCommandQueue(cl_context context, cl_device if (*errcode_ret != CL_SUCCESS) { LOG_ERR(kErrorTemplate, OclError::toString(*errcode_ret), kCreateCommandQueueWithProperties); + + return nullptr; } return result; @@ -366,12 +373,25 @@ cl_int xmrig::OclLib::getProgramInfo(cl_program program, cl_program_info param_n cl_int xmrig::OclLib::releaseCommandQueue(cl_command_queue command_queue) { assert(pReleaseCommandQueue != nullptr); + assert(pGetCommandQueueInfo != nullptr); - const cl_int ret = pReleaseCommandQueue(command_queue); + cl_int ret = pReleaseCommandQueue(command_queue); if (ret != CL_SUCCESS) { LOG_ERR(kErrorTemplate, OclError::toString(ret), kReleaseCommandQueue); } + cl_uint refs = 0; + ret = pGetCommandQueueInfo(command_queue, CL_QUEUE_REFERENCE_COUNT, sizeof(refs), &refs, nullptr); + if (ret == CL_SUCCESS && refs > 0) { + std::this_thread::sleep_for(std::chrono::milliseconds(200)); + } + +# ifndef NDEBUG + ret = pGetCommandQueueInfo(command_queue, CL_QUEUE_REFERENCE_COUNT, sizeof(refs), &refs, nullptr); + assert(ret == CL_SUCCESS); + assert(refs == 0); +# endif + return ret; } @@ -410,6 +430,10 @@ cl_int xmrig::OclLib::releaseMemObject(cl_mem mem_obj) { assert(pReleaseMemObject != nullptr); + if (mem_obj == nullptr) { + return CL_SUCCESS; + } + const cl_int ret = pReleaseMemObject(mem_obj); if (ret != CL_SUCCESS) { LOG_ERR(kErrorTemplate, OclError::toString(ret), kReleaseMemObject); @@ -457,7 +481,15 @@ cl_mem xmrig::OclLib::createBuffer(cl_context context, cl_mem_flags flags, size_ { assert(pCreateBuffer != nullptr); - return pCreateBuffer(context, flags, size, host_ptr, errcode_ret); + auto result = pCreateBuffer(context, flags, size, host_ptr, errcode_ret); + if (*errcode_ret != CL_SUCCESS) { + LOG_ERR(MAGENTA_BG_BOLD(WHITE_BOLD_S " ocl ") RED(" error ") RED_BOLD("%s") RED(" when calling ") RED_BOLD("%s") RED(" with buffer size ") RED_BOLD("%zu"), + OclError::toString(*errcode_ret), kCreateBuffer, size); + + return nullptr; + } + + return result; } diff --git a/src/crypto/cn/CnAlgo.h b/src/crypto/cn/CnAlgo.h index cd4fc007..6df22a86 100644 --- a/src/crypto/cn/CnAlgo.h +++ b/src/crypto/cn/CnAlgo.h @@ -79,6 +79,55 @@ public: return 0; } + inline static uint32_t iterations(Algorithm::Id algo) + { + switch (algo) { + case Algorithm::CN_0: + case Algorithm::CN_1: + case Algorithm::CN_2: + case Algorithm::CN_R: + case Algorithm::CN_WOW: + case Algorithm::CN_RTO: + return CN_ITER; + + case Algorithm::CN_FAST: + case Algorithm::CN_HALF: +# ifdef XMRIG_ALGO_CN_LITE + case Algorithm::CN_LITE_0: + case Algorithm::CN_LITE_1: +# endif +# ifdef XMRIG_ALGO_CN_HEAVY + case Algorithm::CN_HEAVY_0: + case Algorithm::CN_HEAVY_TUBE: + case Algorithm::CN_HEAVY_XHV: +# endif + return CN_ITER / 2; + + case Algorithm::CN_RWZ: + case Algorithm::CN_ZLS: + return 0x60000; + + case Algorithm::CN_XAO: + case Algorithm::CN_DOUBLE: + return CN_ITER * 2; + +# ifdef XMRIG_ALGO_CN_GPU + case Algorithm::CN_GPU: + return 0xC000; +# endif + +# ifdef XMRIG_ALGO_CN_PICO + case Algorithm::CN_PICO_0: + return CN_ITER / 8; + + default: + break; + } +# endif + + return 0; + } + inline static uint32_t mask(Algorithm::Id algo) { # ifdef XMRIG_ALGO_CN_GPU @@ -188,6 +237,7 @@ private: # endif }; + constexpr const static uint32_t m_iterations[] = { CN_ITER, // CN_0 CN_ITER, // CN_1