Added classes IOclRunner, OclBaseRunner, OclCnRunner, OclRxRunner.

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);
 }

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<uint32_t>(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); }

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();
 }
 
 

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;
 };
 

src/backend/opencl/interfaces/IOclRunner.h

@@ -0,0 +1,53 @@
+/* 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_IOCLRUNNER_H
+#define XMRIG_IOCLRUNNER_H
+
+
+#include <stdint.h>
+
+
+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

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)
 

src/backend/opencl/runners/OclBaseRunner.cpp

@@ -0,0 +1,78 @@
+/* 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/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)
+{
+
+}

src/backend/opencl/runners/OclBaseRunner.h

@@ -0,0 +1,68 @@
+/* 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_OCLBASERUNNER_H
+#define XMRIG_OCLBASERUNNER_H
+
+
+#include <string>
+
+
+#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

src/backend/opencl/runners/OclCnRunner.cpp

@@ -0,0 +1,99 @@
+/* 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/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;
+}

src/backend/opencl/runners/OclCnRunner.h

@@ -0,0 +1,57 @@
+/* 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_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

src/backend/opencl/runners/OclRxRunner.cpp

@@ -0,0 +1,36 @@
+/* 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/runners/OclRxRunner.h"
+
+
+xmrig::OclRxRunner::OclRxRunner(size_t index, const OclLaunchData &data) : OclBaseRunner(index, data)
+{
+}
+
+
+bool xmrig::OclRxRunner::selfTest() const
+{
+    return false; // TODO
+}

src/backend/opencl/runners/OclRxRunner.h

@@ -0,0 +1,48 @@
+/* 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_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

src/backend/opencl/wrappers/OclLib.cpp

@@ -23,6 +23,7 @@
  */
 
 
+#include <thread>
 #include <uv.h>
 
 
@@ -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<void**>(&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;
 }
 
 

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