diff --git a/scripts/generate_cl.js b/scripts/generate_cl.js
index 34b33595..5e4c2c4d 100644
--- a/scripts/generate_cl.js
+++ b/scripts/generate_cl.js
@@ -64,8 +64,8 @@ function rx()
         'randomx_jit.cl'
     ]);
 
-    rx = rx.replace(/	#include "fillAes1Rx4.cl"/g, fs.readFileSync('fillAes1Rx4.cl', 'utf8'));
-    rx = rx.replace(/	#include "blake2b_double_block.cl"/g, fs.readFileSync('blake2b_double_block.cl', 'utf8'));
+    rx = rx.replace(/(\t| )*#include "fillAes1Rx4.cl"/g, fs.readFileSync('fillAes1Rx4.cl', 'utf8'));
+    rx = rx.replace(/(\t| )*#include "blake2b_double_block.cl"/g, fs.readFileSync('blake2b_double_block.cl', 'utf8'));
 
     //fs.writeFileSync('randomx_gen.cl', rx);
     fs.writeFileSync('randomx_cl.h', text2h(rx, 'xmrig', 'randomx_cl'));
diff --git a/src/backend/common/Tags.h b/src/backend/common/Tags.h
new file mode 100644
index 00000000..938a1bc4
--- /dev/null
+++ b/src/backend/common/Tags.h
@@ -0,0 +1,44 @@
+/* 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      Lee Clagett <https://github.com/vtnerd>
+ * 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_TAGS_H
+#define XMRIG_TAGS_H
+
+
+namespace xmrig {
+
+
+const char *cpu_tag();
+
+
+#ifdef XMRIG_FEATURE_OPENCL
+const char *ocl_tag();
+#endif
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_TAGS_H */
diff --git a/src/backend/common/Workers.cpp b/src/backend/common/Workers.cpp
index 27a0410c..7ed9bd8e 100644
--- a/src/backend/common/Workers.cpp
+++ b/src/backend/common/Workers.cpp
@@ -151,12 +151,13 @@ xmrig::IWorker *xmrig::Workers<T>::create(Thread<T> *)
 template<class T>
 void xmrig::Workers<T>::onReady(void *arg)
 {
-    Thread<T> *handle = static_cast<Thread<T>* >(arg);
+    auto handle = static_cast<Thread<T>* >(arg);
 
     IWorker *worker = create(handle);
     if (!worker || !worker->selfTest()) {
         LOG_ERR("thread %zu error: \"hash self-test failed\".", worker->id());
 
+        delete worker;
         return;
     }
 
diff --git a/src/backend/common/common.cmake b/src/backend/common/common.cmake
index ddeca0d6..cddd4001 100644
--- a/src/backend/common/common.cmake
+++ b/src/backend/common/common.cmake
@@ -1,5 +1,6 @@
 set(HEADERS_BACKEND_COMMON
     src/backend/common/Hashrate.h
+    src/backend/common/Tags.h
     src/backend/common/interfaces/IBackend.h
     src/backend/common/interfaces/IRxListener.h
     src/backend/common/interfaces/IThread.h
diff --git a/src/backend/cpu/CpuBackend.cpp b/src/backend/cpu/CpuBackend.cpp
index f86a45a8..22ae63d8 100644
--- a/src/backend/cpu/CpuBackend.cpp
+++ b/src/backend/cpu/CpuBackend.cpp
@@ -28,6 +28,7 @@
 
 #include "backend/common/Hashrate.h"
 #include "backend/common/interfaces/IWorker.h"
+#include "backend/common/Tags.h"
 #include "backend/common/Workers.h"
 #include "backend/cpu/Cpu.h"
 #include "backend/cpu/CpuBackend.h"
@@ -196,6 +197,12 @@ public:
 } // namespace xmrig
 
 
+const char *xmrig::cpu_tag()
+{
+    return tag;
+}
+
+
 xmrig::CpuBackend::CpuBackend(Controller *controller) :
     d_ptr(new CpuBackendPrivate(controller))
 {
diff --git a/src/backend/opencl/OclBackend.cpp b/src/backend/opencl/OclBackend.cpp
index 2d066e7c..6471ccd6 100644
--- a/src/backend/opencl/OclBackend.cpp
+++ b/src/backend/opencl/OclBackend.cpp
@@ -28,6 +28,7 @@
 
 #include "backend/common/Hashrate.h"
 #include "backend/common/interfaces/IWorker.h"
+#include "backend/common/Tags.h"
 #include "backend/common/Workers.h"
 #include "backend/opencl/OclBackend.h"
 #include "backend/opencl/OclConfig.h"
@@ -192,6 +193,12 @@ public:
 } // namespace xmrig
 
 
+const char *xmrig::ocl_tag()
+{
+    return tag;
+}
+
+
 xmrig::OclBackend::OclBackend(Controller *controller) :
     d_ptr(new OclBackendPrivate(controller))
 {
diff --git a/src/backend/opencl/OclCache.cpp b/src/backend/opencl/OclCache.cpp
index 876ce74a..7e5504c0 100644
--- a/src/backend/opencl/OclCache.cpp
+++ b/src/backend/opencl/OclCache.cpp
@@ -30,6 +30,7 @@
 
 
 #include "3rdparty/base32/base32.h"
+#include "backend/common/Tags.h"
 #include "backend/opencl/interfaces/IOclRunner.h"
 #include "backend/opencl/OclCache.h"
 #include "backend/opencl/OclLaunchData.h"
@@ -42,13 +43,12 @@
 namespace xmrig {
 
 
-static const char *tag = MAGENTA_BG_BOLD(WHITE_BOLD_S " ocl ");
 static std::mutex mutex;
 
 
 static cl_program createFromSource(const IOclRunner *runner)
 {
-    LOG_INFO("%s GPU " WHITE_BOLD("#%zu") " " YELLOW_BOLD("compiling..."), tag, runner->data().device.index());
+    LOG_INFO("%s GPU " WHITE_BOLD("#%zu") " " YELLOW_BOLD("compiling..."), ocl_tag(), runner->data().device.index());
 
     cl_int ret;
     cl_device_id device = runner->data().device.id();
@@ -68,7 +68,7 @@ static cl_program createFromSource(const IOclRunner *runner)
     }
 
     LOG_INFO("%s GPU " WHITE_BOLD("#%zu") " " GREEN_BOLD("compilation completed") BLACK_BOLD(" (%" PRIu64 " ms)"),
-             tag, runner->data().device.index(), Chrono::steadyMSecs() - ts);
+             ocl_tag(), runner->data().device.index(), Chrono::steadyMSecs() - ts);
 
     return program;
 }
diff --git a/src/backend/opencl/OclCache.h b/src/backend/opencl/OclCache.h
index 75d9e068..e3eea5c3 100644
--- a/src/backend/opencl/OclCache.h
+++ b/src/backend/opencl/OclCache.h
@@ -29,7 +29,7 @@
 #include <string>
 
 
-typedef struct _cl_program *cl_program;
+using cl_program = struct _cl_program *;
 
 
 namespace xmrig {
diff --git a/src/backend/opencl/OclConfig.cpp b/src/backend/opencl/OclConfig.cpp
index bca6e30f..2e27dcd6 100644
--- a/src/backend/opencl/OclConfig.cpp
+++ b/src/backend/opencl/OclConfig.cpp
@@ -166,6 +166,10 @@ std::vector<xmrig::OclLaunchData> xmrig::OclConfig::get(const Miner *miner, cons
             continue;
         }
 
+#       ifdef XMRIG_ALGO_RANDOMX
+        auto dataset = algorithm.family() == Algorithm::RANDOM_X ? std::make_shared<OclRxDataset>() : nullptr;
+#       endif
+
         if (thread.threads().size() > 1) {
             auto interleave = std::make_shared<OclInterleave>(thread.threads().size());
 
@@ -173,11 +177,21 @@ std::vector<xmrig::OclLaunchData> xmrig::OclConfig::get(const Miner *miner, cons
                 OclLaunchData data(miner, algorithm, *this, platform, thread, devices[thread.index()], affinity);
                 data.interleave = interleave;
 
-                out.emplace_back(data);
+#               ifdef XMRIG_ALGO_RANDOMX
+                data.dataset = dataset;
+#               endif
+
+                out.emplace_back(std::move(data));
             }
         }
         else {
-           out.emplace_back(miner, algorithm, *this, platform, thread, devices[thread.index()], thread.threads()[0]);
+            OclLaunchData data(miner, algorithm, *this, platform, thread, devices[thread.index()], thread.threads().front());
+
+#           ifdef XMRIG_ALGO_RANDOMX
+            data.dataset = dataset;
+#           endif
+
+            out.emplace_back(std::move(data));
         }
     }
 
diff --git a/src/backend/opencl/OclLaunchData.h b/src/backend/opencl/OclLaunchData.h
index eb39de10..66c2aaf0 100644
--- a/src/backend/opencl/OclLaunchData.h
+++ b/src/backend/opencl/OclLaunchData.h
@@ -35,6 +35,11 @@
 #include "crypto/common/Nonce.h"
 
 
+#ifdef XMRIG_ALGO_RANDOMX
+#   include "backend/opencl/runners/tools/OclRxDataset.h"
+#endif
+
+
 using cl_context = struct _cl_context *;
 
 
@@ -66,6 +71,10 @@ public:
     const OclPlatform platform;
     const OclThread thread;
     OclInterleavePtr interleave;
+
+#   ifdef XMRIG_ALGO_RANDOMX
+    OclRxDatasetPtr dataset;
+#   endif
 };
 
 
diff --git a/src/backend/opencl/OclThread.cpp b/src/backend/opencl/OclThread.cpp
index 517e98d4..8e4fb844 100644
--- a/src/backend/opencl/OclThread.cpp
+++ b/src/backend/opencl/OclThread.cpp
@@ -61,7 +61,7 @@ xmrig::OclThread::OclThread(const rapidjson::Value &value)
 #   ifdef XMRIG_ALGO_RANDOMX
     m_bfactor     = Json::getUint(value, kBFactor, 6);
     m_gcnAsm      = Json::getUint(value, kGCNAsm, m_gcnAsm);
-    m_datasetHost = Json::getInt(value, kDatasetHost, m_datasetHost);
+    m_datasetHost = Json::getBool(value, kDatasetHost, m_datasetHost);
 #   endif
 
     const rapidjson::Value &si = Json::getArray(value, kStridedIndex);
@@ -134,11 +134,11 @@ rapidjson::Value xmrig::OclThread::toJSON(rapidjson::Document &doc) const
     out.AddMember(StringRef(kUnroll),       unrollFactor(), allocator);
 
 #   ifdef XMRIG_ALGO_RANDOMX
-    if (m_datasetHost != -1) {
-        out.AddMember(StringRef(kBFactor),      bfactor(), allocator);
-        out.AddMember(StringRef(kGCNAsm),       gcnAsm(), allocator);
-        out.AddMember(StringRef(kDatasetHost),  datasetHost(), allocator);
-    }
+//    if (m_datasetHost != -1) {
+//        out.AddMember(StringRef(kBFactor),      bfactor(), allocator);
+//        out.AddMember(StringRef(kGCNAsm),       gcnAsm(), allocator);
+//        out.AddMember(StringRef(kDatasetHost),  isDatasetHost(), allocator);
+//    }
 #   endif
 
     return out;
diff --git a/src/backend/opencl/OclThread.h b/src/backend/opencl/OclThread.h
index 824216da..f219aed2 100644
--- a/src/backend/opencl/OclThread.h
+++ b/src/backend/opencl/OclThread.h
@@ -68,11 +68,11 @@ public:
 
     OclThread(const rapidjson::Value &value);
 
+    inline bool isAsm() const                               { return m_gcnAsm; }
+    inline bool isDatasetHost() const                       { return m_datasetHost; }
     inline bool isValid() const                             { return m_intensity > 0; }
     inline const std::vector<int64_t> &threads() const      { return m_threads; }
     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; }
     inline uint32_t index() const                           { return m_index; }
     inline uint32_t intensity() const                       { return m_intensity; }
     inline uint32_t memChunk() const                        { return m_memChunk; }
@@ -95,11 +95,11 @@ private:
 
     inline void setIntensity(uint32_t intensity)            { m_intensity = intensity / m_worksize * m_worksize; }
 
-    int m_datasetHost               = -1;
+    bool m_datasetHost              = false;
+    bool m_gcnAsm                   = false;
     std::bitset<FIELD_MAX> m_fields = 1;
     std::vector<int64_t> m_threads;
     uint32_t m_bfactor              = 6;
-    uint32_t m_gcnAsm               = 1;
     uint32_t m_index                = 0;
     uint32_t m_intensity            = 0;
     uint32_t m_memChunk             = 2;
diff --git a/src/backend/opencl/OclWorker.cpp b/src/backend/opencl/OclWorker.cpp
index 40545357..4247fb78 100644
--- a/src/backend/opencl/OclWorker.cpp
+++ b/src/backend/opencl/OclWorker.cpp
@@ -26,6 +26,7 @@
 
 #include "backend/opencl/OclWorker.h"
 
+#include "backend/common/Tags.h"
 #include "backend/opencl/runners/OclCnRunner.h"
 #include "base/io/log/Log.h"
 #include "base/tools/Chrono.h"
@@ -35,7 +36,8 @@
 
 
 #ifdef XMRIG_ALGO_RANDOMX
-#   include "backend/opencl/runners/OclRxRunner.h"
+#   include "backend/opencl/runners/OclRxJitRunner.h"
+#   include "backend/opencl/runners/OclRxVmRunner.h"
 #endif
 
 #ifdef XMRIG_ALGO_CN_GPU
@@ -58,6 +60,12 @@ static inline bool isReady()                         { return !Nonce::isPaused()
 static inline uint32_t roundSize(uint32_t intensity) { return kReserveCount / intensity + 1; }
 
 
+static inline void printError(size_t id, const char *error)
+{
+    LOG_ERR("%s" RED_S " thread " RED_BOLD("#%zu") RED_S " failed with error " RED_BOLD("%s"), ocl_tag(), id, error);
+}
+
+
 } // namespace xmrig
 
 
@@ -72,7 +80,12 @@ xmrig::OclWorker::OclWorker(size_t id, const OclLaunchData &data) :
     switch (m_algorithm.family()) {
     case Algorithm::RANDOM_X:
 #       ifdef XMRIG_ALGO_RANDOMX
-        m_runner = new OclRxRunner(id, data);
+        if (data.thread.isAsm() && data.device.vendorId() == OCL_VENDOR_AMD) {
+            m_runner = new OclRxJitRunner(id, data);
+        }
+        else {
+            m_runner = new OclRxVmRunner(id, data);
+        }
 #       endif
         break;
 
@@ -95,9 +108,20 @@ xmrig::OclWorker::OclWorker(size_t id, const OclLaunchData &data) :
         break;
     }
 
-    if (m_runner) {
+    if (!m_runner) {
+        return;
+    }
+
+    try {
+        m_runner->init();
         m_runner->build();
     }
+    catch (std::exception &ex) {
+        printError(id, ex.what());
+
+        delete m_runner;
+        m_runner = nullptr;
+    }
 }
 
 
@@ -109,7 +133,7 @@ xmrig::OclWorker::~OclWorker()
 
 bool xmrig::OclWorker::selfTest()
 {
-    return m_runner && m_runner->selfTest();
+    return m_runner != nullptr;
 }
 
 
@@ -136,7 +160,9 @@ void xmrig::OclWorker::start()
                 m_interleave->resumeDelay(m_id);
             }
 
-            consumeJob();
+            if (!consumeJob()) {
+                return;
+            }
         }
 
         while (!Nonce::isOutdated(Nonce::OPENCL, m_job.sequence())) {
@@ -146,7 +172,12 @@ void xmrig::OclWorker::start()
 
             const uint64_t t = Chrono::steadyMSecs();
 
-            if (!m_runner->run(*m_job.nonce(), results)) {
+            try {
+                m_runner->run(*m_job.nonce(), results);
+            }
+            catch (std::exception &ex) {
+                printError(id(), ex.what());
+
                 return;
             }
 
@@ -160,19 +191,31 @@ void xmrig::OclWorker::start()
             std::this_thread::yield();
         }
 
-        consumeJob();
+        if (!consumeJob()) {
+            return;
+        }
     }
 }
 
 
-void xmrig::OclWorker::consumeJob()
+bool xmrig::OclWorker::consumeJob()
 {
     if (Nonce::sequence(Nonce::OPENCL) == 0) {
-        return;
+        return false;
     }
 
     m_job.add(m_miner->job(), Nonce::sequence(Nonce::OPENCL), roundSize(m_intensity) * m_intensity);
-    m_runner->set(m_job.currentJob(), m_job.blob());
+
+    try {
+        m_runner->set(m_job.currentJob(), m_job.blob());
+    }
+    catch (std::exception &ex) {
+        printError(id(), ex.what());
+
+        return false;
+    }
+
+    return true;
 }
 
 
diff --git a/src/backend/opencl/OclWorker.h b/src/backend/opencl/OclWorker.h
index 95c01d01..e28c784d 100644
--- a/src/backend/opencl/OclWorker.h
+++ b/src/backend/opencl/OclWorker.h
@@ -59,7 +59,7 @@ protected:
     void start() override;
 
 private:
-    void consumeJob();
+    bool consumeJob();
     void storeStats(uint64_t ts);
 
     const Algorithm m_algorithm;
diff --git a/src/backend/opencl/interfaces/IOclRunner.h b/src/backend/opencl/interfaces/IOclRunner.h
index 613244bf..c432ba0a 100644
--- a/src/backend/opencl/interfaces/IOclRunner.h
+++ b/src/backend/opencl/interfaces/IOclRunner.h
@@ -26,10 +26,13 @@
 #define XMRIG_IOCLRUNNER_H
 
 
-#include <stdint.h>
+#include "base/tools/Object.h"
 
 
-typedef struct _cl_context *cl_context;
+#include <cstdint>
+
+
+using cl_context = struct _cl_context *;
 
 
 namespace xmrig {
@@ -43,10 +46,12 @@ class OclLaunchData;
 class IOclRunner
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE(IOclRunner)
+
+    IOclRunner()          = default;
     virtual ~IOclRunner() = default;
 
     virtual bool run(uint32_t nonce, uint32_t *hashOutput)  = 0;
-    virtual bool selfTest() const                           = 0;
     virtual bool set(const Job &job, uint8_t *blob)         = 0;
     virtual cl_context ctx() const                          = 0;
     virtual const Algorithm &algorithm() const              = 0;
@@ -57,9 +62,7 @@ public:
     virtual size_t threadId() const                         = 0;
     virtual uint32_t deviceIndex() const                    = 0;
     virtual void build()                                    = 0;
-
-protected:
-    virtual bool isReadyToBuild() const                     = 0;
+    virtual void init()                                     = 0;
 };
 
 
diff --git a/src/backend/opencl/kernels/Cn00RyoKernel.cpp b/src/backend/opencl/kernels/Cn00RyoKernel.cpp
index 31b33066..df987fa4 100644
--- a/src/backend/opencl/kernels/Cn00RyoKernel.cpp
+++ b/src/backend/opencl/kernels/Cn00RyoKernel.cpp
@@ -27,17 +27,18 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
-bool xmrig::Cn00RyoKernel::enqueue(cl_command_queue queue, size_t threads)
+void xmrig::Cn00RyoKernel::enqueue(cl_command_queue queue, size_t threads)
 {
     const size_t gthreads = threads * 64;
     const size_t lthreads = 64;
 
-    return enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
 }
 
 
 // __kernel void cn00(__global int *Scratchpad, __global ulong *states)
-bool xmrig::Cn00RyoKernel::setArgs(cl_mem scratchpads, cl_mem states)
+void xmrig::Cn00RyoKernel::setArgs(cl_mem scratchpads, cl_mem states)
 {
-    return setArg(0, sizeof(cl_mem), &scratchpads) && setArg(1, sizeof(cl_mem), &states);
+    setArg(0, sizeof(cl_mem), &scratchpads);
+    setArg(1, sizeof(cl_mem), &states);
 }
diff --git a/src/backend/opencl/kernels/Cn00RyoKernel.h b/src/backend/opencl/kernels/Cn00RyoKernel.h
index 2ad35e35..366f644e 100644
--- a/src/backend/opencl/kernels/Cn00RyoKernel.h
+++ b/src/backend/opencl/kernels/Cn00RyoKernel.h
@@ -37,8 +37,8 @@ class Cn00RyoKernel : public OclKernel
 public:
     inline Cn00RyoKernel(cl_program program) : OclKernel(program, "cn00") {}
 
-    bool enqueue(cl_command_queue queue, size_t threads);
-    bool setArgs(cl_mem scratchpads, cl_mem states);
+    void enqueue(cl_command_queue queue, size_t threads);
+    void setArgs(cl_mem scratchpads, cl_mem states);
 };
 
 
diff --git a/src/backend/opencl/kernels/Cn0Kernel.cpp b/src/backend/opencl/kernels/Cn0Kernel.cpp
index 41040324..a93e1005 100644
--- a/src/backend/opencl/kernels/Cn0Kernel.cpp
+++ b/src/backend/opencl/kernels/Cn0Kernel.cpp
@@ -27,21 +27,21 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
-bool xmrig::Cn0Kernel::enqueue(cl_command_queue queue, uint32_t nonce, size_t threads)
+void xmrig::Cn0Kernel::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);
+    enqueueNDRange(queue, 2, offset, gthreads, lthreads);
 }
 
 
 // __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)
+void 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);
+    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
index 91328d17..1bb9a37a 100644
--- a/src/backend/opencl/kernels/Cn0Kernel.h
+++ b/src/backend/opencl/kernels/Cn0Kernel.h
@@ -37,8 +37,8 @@ class Cn0Kernel : public OclKernel
 public:
     inline Cn0Kernel(cl_program program) : OclKernel(program, "cn0") {}
 
-    bool enqueue(cl_command_queue queue, uint32_t nonce, size_t threads);
-    bool setArgs(cl_mem input, cl_mem scratchpads, cl_mem states, uint32_t threads);
+    void enqueue(cl_command_queue queue, uint32_t nonce, size_t threads);
+    void setArgs(cl_mem input, cl_mem scratchpads, cl_mem states, uint32_t threads);
 };
 
 
diff --git a/src/backend/opencl/kernels/Cn1Kernel.cpp b/src/backend/opencl/kernels/Cn1Kernel.cpp
index 375e7f2d..fad7889c 100644
--- a/src/backend/opencl/kernels/Cn1Kernel.cpp
+++ b/src/backend/opencl/kernels/Cn1Kernel.cpp
@@ -43,21 +43,21 @@ xmrig::Cn1Kernel::Cn1Kernel(cl_program program, uint64_t height)
 }
 
 
-bool xmrig::Cn1Kernel::enqueue(cl_command_queue queue, uint32_t nonce, size_t threads, size_t worksize)
+void xmrig::Cn1Kernel::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, &gthreads, &lthreads);
+    enqueueNDRange(queue, 1, &offset, &gthreads, &lthreads);
 }
 
 
 // __kernel void cn1(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states, uint Threads)
-bool xmrig::Cn1Kernel::setArgs(cl_mem input, cl_mem scratchpads, cl_mem states, uint32_t threads)
+void xmrig::Cn1Kernel::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);
+    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/Cn1Kernel.h b/src/backend/opencl/kernels/Cn1Kernel.h
index 523ae604..08d33de2 100644
--- a/src/backend/opencl/kernels/Cn1Kernel.h
+++ b/src/backend/opencl/kernels/Cn1Kernel.h
@@ -38,8 +38,8 @@ public:
     Cn1Kernel(cl_program program);
     Cn1Kernel(cl_program program, uint64_t height);
 
-    bool enqueue(cl_command_queue queue, uint32_t nonce, size_t threads, size_t worksize);
-    bool setArgs(cl_mem input, cl_mem scratchpads, cl_mem states, uint32_t threads);
+    void enqueue(cl_command_queue queue, uint32_t nonce, size_t threads, size_t worksize);
+    void setArgs(cl_mem input, cl_mem scratchpads, cl_mem states, uint32_t threads);
 };
 
 
diff --git a/src/backend/opencl/kernels/Cn1RyoKernel.cpp b/src/backend/opencl/kernels/Cn1RyoKernel.cpp
index bec96b7c..50254543 100644
--- a/src/backend/opencl/kernels/Cn1RyoKernel.cpp
+++ b/src/backend/opencl/kernels/Cn1RyoKernel.cpp
@@ -30,19 +30,19 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
-bool xmrig::Cn1RyoKernel::enqueue(cl_command_queue queue, size_t threads, size_t worksize)
+void xmrig::Cn1RyoKernel::enqueue(cl_command_queue queue, size_t threads, size_t worksize)
 {
     const size_t gthreads = threads * 16;
     const size_t lthreads = worksize * 16;
 
-    return enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
 }
 
 
 // __kernel void cn1(__global int *lpad_in, __global int *spad, uint numThreads)
-bool xmrig::Cn1RyoKernel::setArgs(cl_mem scratchpads, cl_mem states, uint32_t threads)
+void xmrig::Cn1RyoKernel::setArgs(cl_mem scratchpads, cl_mem states, uint32_t threads)
 {
-    return setArg(0, sizeof(cl_mem), &scratchpads) &&
-           setArg(1, sizeof(cl_mem), &states) &&
-           setArg(2, sizeof(uint32_t), &threads);
+    setArg(0, sizeof(cl_mem), &scratchpads);
+    setArg(1, sizeof(cl_mem), &states);
+    setArg(2, sizeof(uint32_t), &threads);
 }
diff --git a/src/backend/opencl/kernels/Cn1RyoKernel.h b/src/backend/opencl/kernels/Cn1RyoKernel.h
index 6350f516..31714f1e 100644
--- a/src/backend/opencl/kernels/Cn1RyoKernel.h
+++ b/src/backend/opencl/kernels/Cn1RyoKernel.h
@@ -37,8 +37,8 @@ class Cn1RyoKernel : public OclKernel
 public:
     inline Cn1RyoKernel(cl_program program) : OclKernel(program, "cn1") {}
 
-    bool enqueue(cl_command_queue queue, size_t threads, size_t worksize);
-    bool setArgs(cl_mem scratchpads, cl_mem states, uint32_t threads);
+    void enqueue(cl_command_queue queue, size_t threads, size_t worksize);
+    void setArgs(cl_mem scratchpads, cl_mem states, uint32_t threads);
 };
 
 
diff --git a/src/backend/opencl/kernels/Cn2Kernel.cpp b/src/backend/opencl/kernels/Cn2Kernel.cpp
index eaa3b7db..1dbb03b1 100644
--- a/src/backend/opencl/kernels/Cn2Kernel.cpp
+++ b/src/backend/opencl/kernels/Cn2Kernel.cpp
@@ -27,28 +27,24 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
-bool xmrig::Cn2Kernel::enqueue(cl_command_queue queue, uint32_t nonce, size_t threads)
+void 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);
+    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<cl_mem> &branches, uint32_t threads)
+void xmrig::Cn2Kernel::setArgs(cl_mem scratchpads, cl_mem states, const std::vector<cl_mem> &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;
-    }
+    setArg(0, sizeof(cl_mem), &scratchpads);
+    setArg(1, sizeof(cl_mem), &states);
+    setArg(6, sizeof(uint32_t), &threads);
 
     for (uint32_t i = 0; i < branches.size(); ++i) {
-        if (!setArg(i + 2, sizeof(cl_mem), &branches[i])) {
-            return false;
-        }
+        setArg(i + 2, sizeof(cl_mem), &branches[i]);
     }
-
-    return true;
 }
diff --git a/src/backend/opencl/kernels/Cn2Kernel.h b/src/backend/opencl/kernels/Cn2Kernel.h
index d409b611..cb4aaede 100644
--- a/src/backend/opencl/kernels/Cn2Kernel.h
+++ b/src/backend/opencl/kernels/Cn2Kernel.h
@@ -37,8 +37,8 @@ class Cn2Kernel : public OclKernel
 public:
     inline Cn2Kernel(cl_program program) : OclKernel(program, "cn2") {}
 
-    bool enqueue(cl_command_queue queue, uint32_t nonce, size_t threads);
-    bool setArgs(cl_mem scratchpads, cl_mem states, const std::vector<cl_mem> &branches, uint32_t threads);
+    void enqueue(cl_command_queue queue, uint32_t nonce, size_t threads);
+    void setArgs(cl_mem scratchpads, cl_mem states, const std::vector<cl_mem> &branches, uint32_t threads);
 };
 
 
diff --git a/src/backend/opencl/kernels/Cn2RyoKernel.cpp b/src/backend/opencl/kernels/Cn2RyoKernel.cpp
index 8824ec26..e294eb24 100644
--- a/src/backend/opencl/kernels/Cn2RyoKernel.cpp
+++ b/src/backend/opencl/kernels/Cn2RyoKernel.cpp
@@ -27,22 +27,27 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
-bool xmrig::Cn2RyoKernel::enqueue(cl_command_queue queue, uint32_t nonce, size_t threads)
+void xmrig::Cn2RyoKernel::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);
+    enqueueNDRange(queue, 2, offset, gthreads, lthreads);
 }
 
 
 // __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global uint *output, ulong Target, uint Threads)
-bool xmrig::Cn2RyoKernel::setArgs(cl_mem scratchpads, cl_mem states, cl_mem output, uint64_t target, uint32_t threads)
+void xmrig::Cn2RyoKernel::setArgs(cl_mem scratchpads, cl_mem states, cl_mem output, uint32_t threads)
 {
-    return setArg(0, sizeof(cl_mem), &scratchpads) &&
-           setArg(1, sizeof(cl_mem), &states) &&
-           setArg(2, sizeof(cl_mem), &output) &&
-           setArg(3, sizeof(cl_ulong), &target) &&
-           setArg(4, sizeof(uint32_t), &threads);
+    setArg(0, sizeof(cl_mem), &scratchpads);
+    setArg(1, sizeof(cl_mem), &states);
+    setArg(2, sizeof(cl_mem), &output);
+    setArg(4, sizeof(uint32_t), &threads);
+}
+
+
+void xmrig::Cn2RyoKernel::setTarget(uint64_t target)
+{
+    setArg(3, sizeof(cl_ulong), &target);
 }
diff --git a/src/backend/opencl/kernels/Cn2RyoKernel.h b/src/backend/opencl/kernels/Cn2RyoKernel.h
index 4dbf2a9c..2ef85bcb 100644
--- a/src/backend/opencl/kernels/Cn2RyoKernel.h
+++ b/src/backend/opencl/kernels/Cn2RyoKernel.h
@@ -37,8 +37,9 @@ class Cn2RyoKernel : public OclKernel
 public:
     inline Cn2RyoKernel(cl_program program) : OclKernel(program, "cn2") {}
 
-    bool enqueue(cl_command_queue queue, uint32_t nonce, size_t threads);
-    bool setArgs(cl_mem scratchpads, cl_mem states, cl_mem output, uint64_t target, uint32_t threads);
+    void enqueue(cl_command_queue queue, uint32_t nonce, size_t threads);
+    void setArgs(cl_mem scratchpads, cl_mem states, cl_mem output, uint32_t threads);
+    void setTarget(uint64_t target);
 };
 
 
diff --git a/src/backend/opencl/kernels/CnBranchKernel.cpp b/src/backend/opencl/kernels/CnBranchKernel.cpp
index 04dd830e..a0813a74 100644
--- a/src/backend/opencl/kernels/CnBranchKernel.cpp
+++ b/src/backend/opencl/kernels/CnBranchKernel.cpp
@@ -41,22 +41,27 @@ xmrig::CnBranchKernel::CnBranchKernel(size_t index, cl_program program) : OclKer
 }
 
 
-bool xmrig::CnBranchKernel::enqueue(cl_command_queue queue, uint32_t nonce, size_t threads, size_t worksize)
+void 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, &gthreads, &lthreads);
+    enqueueNDRange(queue, 1, &offset, &gthreads, &lthreads);
 }
 
 
 // __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)
+void xmrig::CnBranchKernel::setArgs(cl_mem states, cl_mem branch, cl_mem output, 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);
+    setArg(0, sizeof(cl_mem), &states);
+    setArg(1, sizeof(cl_mem), &branch);
+    setArg(2, sizeof(cl_mem), &output);
+    setArg(4, sizeof(cl_uint), &threads);
+}
+
+
+void xmrig::CnBranchKernel::setTarget(uint64_t target)
+{
+    setArg(3, sizeof(cl_ulong), &target);
 }
diff --git a/src/backend/opencl/kernels/CnBranchKernel.h b/src/backend/opencl/kernels/CnBranchKernel.h
index 496d7d2d..a52c928d 100644
--- a/src/backend/opencl/kernels/CnBranchKernel.h
+++ b/src/backend/opencl/kernels/CnBranchKernel.h
@@ -36,8 +36,9 @@ 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);
+    void enqueue(cl_command_queue queue, uint32_t nonce, size_t threads, size_t worksize);
+    void setArgs(cl_mem states, cl_mem branch, cl_mem output, uint32_t threads);
+    void setTarget(uint64_t target);
 };
 
 
diff --git a/src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.cpp b/src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.cpp
new file mode 100644
index 00000000..aea009e5
--- /dev/null
+++ b/src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.cpp
@@ -0,0 +1,37 @@
+/* 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/Blake2bHashRegistersKernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+// __kernel void blake2b_hash_registers_32(__global void *out, __global const void* in, uint inStrideBytes)
+// __kernel void blake2b_hash_registers_64(__global void *out, __global const void* in, uint inStrideBytes)
+void xmrig::Blake2bHashRegistersKernel::setArgs(cl_mem out, cl_mem in, uint32_t inStrideBytes)
+{
+    setArg(0, sizeof(cl_mem), &out);
+    setArg(1, sizeof(cl_mem), &in);
+    setArg(2, sizeof(uint32_t), &inStrideBytes);
+}
diff --git a/src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h b/src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h
new file mode 100644
index 00000000..f1f316c4
--- /dev/null
+++ b/src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h
@@ -0,0 +1,47 @@
+/* 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_BLAKE2BHASHREGISTERSKERNEL_H
+#define XMRIG_BLAKE2BHASHREGISTERSKERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+class Blake2bHashRegistersKernel : public OclKernel
+{
+public:
+    inline Blake2bHashRegistersKernel(cl_program program, const char *name) : OclKernel(program, name) {}
+
+    void setArgs(cl_mem out, cl_mem in, uint32_t inStrideBytes);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_BLAKE2BHASHREGISTERSKERNEL_H */
diff --git a/src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.cpp b/src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.cpp
new file mode 100644
index 00000000..334da909
--- /dev/null
+++ b/src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.cpp
@@ -0,0 +1,35 @@
+/* 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/Blake2bInitialHashKernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+// __kernel void blake2b_initial_hash(__global void *out, __global const void* blockTemplate, uint blockTemplateSize, uint start_nonce)
+void xmrig::Blake2bInitialHashKernel::setArgs(cl_mem out, cl_mem blockTemplate)
+{
+    setArg(0, sizeof(cl_mem), &out);
+    setArg(1, sizeof(cl_mem), &blockTemplate);
+}
diff --git a/src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.h b/src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.h
new file mode 100644
index 00000000..0eba796b
--- /dev/null
+++ b/src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.h
@@ -0,0 +1,47 @@
+/* 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_BLAKE2BINITIALHASHKERNEL_H
+#define XMRIG_BLAKE2BINITIALHASHKERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+class Blake2bInitialHashKernel : public OclKernel
+{
+public:
+    inline Blake2bInitialHashKernel(cl_program program) : OclKernel(program, "blake2b_initial_hash") {}
+
+    void setArgs(cl_mem out, cl_mem blockTemplate);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_BLAKE2BINITIALHASHKERNEL_H */
diff --git a/src/backend/opencl/kernels/rx/ExecuteVmKernel.cpp b/src/backend/opencl/kernels/rx/ExecuteVmKernel.cpp
new file mode 100644
index 00000000..c319757b
--- /dev/null
+++ b/src/backend/opencl/kernels/rx/ExecuteVmKernel.cpp
@@ -0,0 +1,38 @@
+/* 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/ExecuteVmKernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+// __kernel void execute_vm(__global void* vm_states, __global void* rounding, __global void* scratchpads, __global const void* dataset_ptr, uint32_t batch_size, uint32_t num_iterations, uint32_t first, uint32_t last)
+void xmrig::ExecuteVmKernel::setArgs(cl_mem vm_states, cl_mem rounding, cl_mem scratchpads, cl_mem dataset_ptr, uint32_t batch_size)
+{
+    setArg(0, sizeof(cl_mem), &vm_states);
+    setArg(1, sizeof(cl_mem), &rounding);
+    setArg(2, sizeof(cl_mem), &scratchpads);
+    setArg(3, sizeof(cl_mem), &dataset_ptr);
+    setArg(4, sizeof(uint32_t), &batch_size);
+}
diff --git a/src/backend/opencl/kernels/rx/ExecuteVmKernel.h b/src/backend/opencl/kernels/rx/ExecuteVmKernel.h
new file mode 100644
index 00000000..3f9215ef
--- /dev/null
+++ b/src/backend/opencl/kernels/rx/ExecuteVmKernel.h
@@ -0,0 +1,47 @@
+/* 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_EXECUTEVMKERNEL_H
+#define XMRIG_EXECUTEVMKERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+class ExecuteVmKernel : public OclKernel
+{
+public:
+    inline ExecuteVmKernel(cl_program program) : OclKernel(program, "execute_vm") {}
+
+    void setArgs(cl_mem vm_states, cl_mem rounding, cl_mem scratchpads, cl_mem dataset_ptr, uint32_t batch_size);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_EXECUTEVMKERNEL_H */
diff --git a/src/backend/opencl/kernels/rx/FillAesKernel.cpp b/src/backend/opencl/kernels/rx/FillAesKernel.cpp
new file mode 100644
index 00000000..ec960bfa
--- /dev/null
+++ b/src/backend/opencl/kernels/rx/FillAesKernel.cpp
@@ -0,0 +1,38 @@
+/* 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/FillAesKernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+// __kernel void fillAes1Rx4_scratchpad(__global void* state, __global void* out, uint batch_size, uint rx_version)
+// __kernel void fillAes4Rx4_entropy(__global void* state, __global void* out, uint batch_size, uint rx_version)
+void xmrig::FillAesKernel::setArgs(cl_mem state, cl_mem out, uint32_t batch_size, uint32_t rx_version)
+{
+    setArg(0, sizeof(cl_mem), &state);
+    setArg(1, sizeof(cl_mem), &out);
+    setArg(2, sizeof(uint32_t), &batch_size);
+    setArg(3, sizeof(uint32_t), &rx_version);
+}
diff --git a/src/backend/opencl/kernels/rx/FillAesKernel.h b/src/backend/opencl/kernels/rx/FillAesKernel.h
new file mode 100644
index 00000000..8d84a706
--- /dev/null
+++ b/src/backend/opencl/kernels/rx/FillAesKernel.h
@@ -0,0 +1,47 @@
+/* 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_FILLAESKERNEL_H
+#define XMRIG_FILLAESKERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+class FillAesKernel : public OclKernel
+{
+public:
+    inline FillAesKernel(cl_program program, const char *name) : OclKernel(program, name) {}
+
+    void setArgs(cl_mem state, cl_mem out, uint32_t batch_size, uint32_t rx_version);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_FILLAESKERNEL_H */
diff --git a/src/backend/opencl/kernels/rx/FindSharesKernel.cpp b/src/backend/opencl/kernels/rx/FindSharesKernel.cpp
new file mode 100644
index 00000000..fa50619e
--- /dev/null
+++ b/src/backend/opencl/kernels/rx/FindSharesKernel.cpp
@@ -0,0 +1,35 @@
+/* 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/FindSharesKernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+// __kernel void find_shares(__global const uint64_t* hashes, uint64_t target, uint32_t start_nonce, __global uint32_t* shares)
+void xmrig::FindSharesKernel::setArgs(cl_mem hashes, cl_mem shares)
+{
+    setArg(0, sizeof(cl_mem), &hashes);
+    setArg(3, sizeof(cl_mem), &shares);
+}
diff --git a/src/backend/opencl/kernels/rx/FindSharesKernel.h b/src/backend/opencl/kernels/rx/FindSharesKernel.h
new file mode 100644
index 00000000..3c10284e
--- /dev/null
+++ b/src/backend/opencl/kernels/rx/FindSharesKernel.h
@@ -0,0 +1,47 @@
+/* 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_FINDSHARESKERNEL_H
+#define XMRIG_FINDSHARESKERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+class FindSharesKernel : public OclKernel
+{
+public:
+    inline FindSharesKernel(cl_program program) : OclKernel(program, "find_shares") {}
+
+    void setArgs(cl_mem hashes, cl_mem shares);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_FINDSHARESKERNEL_H */
diff --git a/src/backend/opencl/kernels/rx/HashAesKernel.cpp b/src/backend/opencl/kernels/rx/HashAesKernel.cpp
new file mode 100644
index 00000000..3a512269
--- /dev/null
+++ b/src/backend/opencl/kernels/rx/HashAesKernel.cpp
@@ -0,0 +1,40 @@
+/* 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/HashAesKernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+// __kernel void hashAes1Rx4(__global const void* input, __global void* hash, uint hashOffsetBytes, uint hashStrideBytes, uint batch_size)
+void xmrig::HashAesKernel::setArgs(cl_mem input, cl_mem hash, uint32_t hashStrideBytes, uint32_t batch_size)
+{
+    const uint32_t hashOffsetBytes = 192;
+
+    setArg(0, sizeof(cl_mem), &input);
+    setArg(1, sizeof(cl_mem), &hash);
+    setArg(2, sizeof(uint32_t), &hashOffsetBytes);
+    setArg(3, sizeof(uint32_t), &hashStrideBytes);
+    setArg(4, sizeof(uint32_t), &batch_size);
+}
diff --git a/src/backend/opencl/kernels/rx/HashAesKernel.h b/src/backend/opencl/kernels/rx/HashAesKernel.h
new file mode 100644
index 00000000..33d9f17d
--- /dev/null
+++ b/src/backend/opencl/kernels/rx/HashAesKernel.h
@@ -0,0 +1,47 @@
+/* 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_HASHAESKERNEL_H
+#define XMRIG_HASHAESKERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+class HashAesKernel : public OclKernel
+{
+public:
+    inline HashAesKernel(cl_program program) : OclKernel(program, "hashAes1Rx4") {}
+
+    void setArgs(cl_mem input, cl_mem hash, uint32_t hashStrideBytes, uint32_t batch_size);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_HASHAESKERNEL_H */
diff --git a/src/backend/opencl/kernels/rx/InitVmKernel.cpp b/src/backend/opencl/kernels/rx/InitVmKernel.cpp
new file mode 100644
index 00000000..2c083fd5
--- /dev/null
+++ b/src/backend/opencl/kernels/rx/InitVmKernel.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/kernels/rx/InitVmKernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+// __kernel void init_vm(__global const void* entropy_data, __global void* vm_states, __global uint32_t* rounding, uint32_t iteration)
+void xmrig::InitVmKernel::setArgs(cl_mem entropy_data, cl_mem vm_states, cl_mem rounding)
+{
+    setArg(0, sizeof(cl_mem), &entropy_data);
+    setArg(1, sizeof(cl_mem), &vm_states);
+    setArg(2, sizeof(cl_mem), &rounding);
+}
diff --git a/src/backend/opencl/kernels/rx/InitVmKernel.h b/src/backend/opencl/kernels/rx/InitVmKernel.h
new file mode 100644
index 00000000..b1fd6f8a
--- /dev/null
+++ b/src/backend/opencl/kernels/rx/InitVmKernel.h
@@ -0,0 +1,47 @@
+/* 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_INITVMKERNEL_H
+#define XMRIG_INITVMKERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+class InitVmKernel : public OclKernel
+{
+public:
+    inline InitVmKernel(cl_program program) : OclKernel(program, "init_vm") {}
+
+    void setArgs(cl_mem entropy_data, cl_mem vm_states, cl_mem rounding);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_INITVMKERNEL_H */
diff --git a/src/backend/opencl/opencl.cmake b/src/backend/opencl/opencl.cmake
index 3a8c9596..cd0682d9 100644
--- a/src/backend/opencl/opencl.cmake
+++ b/src/backend/opencl/opencl.cmake
@@ -65,8 +65,33 @@ if (WITH_OPENCL)
    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)
+       list(APPEND HEADERS_BACKEND_OPENCL
+           src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h
+           src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.h
+           src/backend/opencl/kernels/rx/ExecuteVmKernel.h
+           src/backend/opencl/kernels/rx/FillAesKernel.h
+           src/backend/opencl/kernels/rx/FindSharesKernel.h
+           src/backend/opencl/kernels/rx/HashAesKernel.cpp
+           src/backend/opencl/kernels/rx/InitVmKernel.h
+           src/backend/opencl/runners/OclRxBaseRunner.h
+           src/backend/opencl/runners/OclRxJitRunner.h
+           src/backend/opencl/runners/OclRxVmRunner.h
+           src/backend/opencl/runners/tools/OclRxDataset.h
+           )
+
+       list(APPEND SOURCES_BACKEND_OPENCL
+           src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.cpp
+           src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.cpp
+           src/backend/opencl/kernels/rx/ExecuteVmKernel.cpp
+           src/backend/opencl/kernels/rx/FillAesKernel.cpp
+           src/backend/opencl/kernels/rx/FindSharesKernel.cpp
+           src/backend/opencl/kernels/rx/HashAesKernel.cpp
+           src/backend/opencl/kernels/rx/InitVmKernel.cpp
+           src/backend/opencl/runners/OclRxBaseRunner.cpp
+           src/backend/opencl/runners/OclRxJitRunner.cpp
+           src/backend/opencl/runners/OclRxVmRunner.cpp
+           src/backend/opencl/runners/tools/OclRxDataset.cpp
+           )
    endif()
 
    if (WITH_CN_GPU AND CMAKE_SIZEOF_VOID_P EQUAL 8)
diff --git a/src/backend/opencl/runners/OclBaseRunner.cpp b/src/backend/opencl/runners/OclBaseRunner.cpp
index ccbd4c48..94c6a68c 100644
--- a/src/backend/opencl/runners/OclBaseRunner.cpp
+++ b/src/backend/opencl/runners/OclBaseRunner.cpp
@@ -30,6 +30,7 @@
 #include "backend/opencl/wrappers/OclLib.h"
 #include "base/io/log/Log.h"
 #include "base/net/stratum/Job.h"
+#include "backend/opencl/wrappers/OclError.h"
 
 
 xmrig::OclBaseRunner::OclBaseRunner(size_t id, const OclLaunchData &data) :
@@ -39,15 +40,6 @@ xmrig::OclBaseRunner::OclBaseRunner(size_t id, const OclLaunchData &data) :
     m_data(data),
     m_threadId(id)
 {
-    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);
-
     m_deviceKey = data.device.name();
 
 #   ifdef XMRIG_STRICT_OPENCL_CACHE
@@ -73,18 +65,6 @@ xmrig::OclBaseRunner::~OclBaseRunner()
 }
 
 
-bool xmrig::OclBaseRunner::isReadyToBuild() const
-{
-    return m_queue != nullptr && m_input != nullptr && m_output != nullptr && !m_options.empty() && m_source != nullptr;
-}
-
-
-bool xmrig::OclBaseRunner::selfTest() const
-{
-    return isReadyToBuild() && m_program != nullptr;
-}
-
-
 uint32_t xmrig::OclBaseRunner::deviceIndex() const
 {
     return data().thread.index();
@@ -93,9 +73,35 @@ uint32_t xmrig::OclBaseRunner::deviceIndex() const
 
 void xmrig::OclBaseRunner::build()
 {
-    if (!isReadyToBuild()) {
-        return;
-    }
-
     m_program = OclCache::build(this);
+
+    if (m_program == nullptr) {
+        throw std::runtime_error(OclError::toString(CL_INVALID_PROGRAM));
+    }
+}
+
+
+void xmrig::OclBaseRunner::init()
+{
+    m_queue  = OclLib::createCommandQueue(m_ctx, data().device.id());
+    m_input  = OclLib::createBuffer(m_ctx, CL_MEM_READ_ONLY, Job::kMaxBlobSize);
+    m_output = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, sizeof(cl_uint) * 0x100);
+}
+
+
+void xmrig::OclBaseRunner::enqueueReadBuffer(cl_mem buffer, cl_bool blocking_read, size_t offset, size_t size, void *ptr)
+{
+    const cl_int ret = OclLib::enqueueReadBuffer(m_queue, buffer, blocking_read, offset, size, ptr, 0, nullptr, nullptr);
+    if (ret != CL_SUCCESS) {
+        throw std::runtime_error(OclError::toString(ret));
+    }
+}
+
+
+void xmrig::OclBaseRunner::enqueueWriteBuffer(cl_mem buffer, cl_bool blocking_write, size_t offset, size_t size, const void *ptr)
+{
+    const cl_int ret = OclLib::enqueueWriteBuffer(m_queue, buffer, blocking_write, offset, size, ptr, 0, nullptr, nullptr);
+    if (ret != CL_SUCCESS) {
+        throw std::runtime_error(OclError::toString(ret));
+    }
 }
diff --git a/src/backend/opencl/runners/OclBaseRunner.h b/src/backend/opencl/runners/OclBaseRunner.h
index 95f0f778..ee3d4b5b 100644
--- a/src/backend/opencl/runners/OclBaseRunner.h
+++ b/src/backend/opencl/runners/OclBaseRunner.h
@@ -43,16 +43,11 @@ class OclLaunchData;
 class OclBaseRunner : public IOclRunner
 {
 public:
-    OclBaseRunner()                           = delete;
-    OclBaseRunner(const OclBaseRunner &other) = delete;
-    OclBaseRunner(OclBaseRunner &&other)      = delete;
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(OclBaseRunner)
+
     OclBaseRunner(size_t id, const OclLaunchData &data);
-
     ~OclBaseRunner() override;
 
-    OclBaseRunner &operator=(const OclBaseRunner &other) = delete;
-    OclBaseRunner &operator=(OclBaseRunner &&other)      = delete;
-
 protected:
     inline cl_context ctx() const override              { return m_ctx; }
     inline const Algorithm &algorithm() const override  { return m_algorithm; }
@@ -62,12 +57,14 @@ protected:
     inline const OclLaunchData &data() const override   { return m_data; }
     inline size_t threadId() const override             { return m_threadId; }
 
-    bool isReadyToBuild() const override;
-    bool selfTest() const override;
     uint32_t deviceIndex() const override;
     void build() override;
+    void init() override;
 
 protected:
+    void enqueueReadBuffer(cl_mem buffer, cl_bool blocking_read, size_t offset, size_t size, void *ptr);
+    void enqueueWriteBuffer(cl_mem buffer, cl_bool blocking_write, size_t offset, size_t size, const void *ptr);
+
     Algorithm m_algorithm;
     cl_command_queue m_queue    = nullptr;
     cl_context m_ctx;
diff --git a/src/backend/opencl/runners/OclCnRunner.cpp b/src/backend/opencl/runners/OclCnRunner.cpp
index 30cb275c..49c60099 100644
--- a/src/backend/opencl/runners/OclCnRunner.cpp
+++ b/src/backend/opencl/runners/OclCnRunner.cpp
@@ -39,30 +39,6 @@
 
 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);
-    if (ret != CL_SUCCESS) {
-        return;
-    }
-
-    for (size_t i = 0; i < BRANCH_MAX; ++i) {
-        m_branches[i] = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, sizeof(cl_uint) * (g_thd + 2), nullptr, &ret);
-        if (ret != CL_SUCCESS) {
-            return;
-        }
-    }
-
     uint32_t stridedIndex = data.thread.stridedIndex();
     if (data.device.vendorId() == OCL_VENDOR_NVIDIA) {
         stridedIndex = 0;
@@ -105,18 +81,6 @@ xmrig::OclCnRunner::~OclCnRunner()
 }
 
 
-bool xmrig::OclCnRunner::isReadyToBuild() const
-{
-    return OclBaseRunner::isReadyToBuild() &&
-            m_scratchpads                   != nullptr &&
-            m_states                        != nullptr &&
-            m_branches[BRANCH_BLAKE_256]    != nullptr &&
-            m_branches[BRANCH_GROESTL_256]  != nullptr &&
-            m_branches[BRANCH_JH_256]       != nullptr &&
-            m_branches[BRANCH_SKEIN_512]    != nullptr;
-}
-
-
 bool xmrig::OclCnRunner::run(uint32_t nonce, uint32_t *hashOutput)
 {
     static const cl_uint zero = 0;
@@ -128,36 +92,20 @@ bool xmrig::OclCnRunner::run(uint32_t nonce, uint32_t *hashOutput)
     assert(g_thd % w_size == 0);
 
     for (size_t i = 0; i < BRANCH_MAX; ++i) {
-        if (OclLib::enqueueWriteBuffer(m_queue, m_branches[i], CL_FALSE, sizeof(cl_uint) * g_intensity, sizeof(cl_uint), &zero, 0, nullptr, nullptr) != CL_SUCCESS) {
-            return false;
-        }
+        enqueueWriteBuffer(m_branches[i], CL_FALSE, sizeof(cl_uint) * g_intensity, sizeof(cl_uint), &zero);
     }
 
-    if (OclLib::enqueueWriteBuffer(m_queue, m_output, CL_FALSE, sizeof(cl_uint) * 0xFF, sizeof(cl_uint), &zero, 0, nullptr, nullptr) != CL_SUCCESS) {
-        return false;
+    enqueueWriteBuffer(m_output, CL_FALSE, sizeof(cl_uint) * 0xFF, sizeof(cl_uint), &zero);
+
+    m_cn0->enqueue(m_queue, nonce, g_thd);
+    m_cn1->enqueue(m_queue, nonce, g_thd, w_size);
+    m_cn2->enqueue(m_queue, nonce, g_thd);
+
+    for (auto kernel : m_branchKernels) {
+        kernel->enqueue(m_queue, nonce, g_thd, w_size);
     }
 
-    if (!m_cn0->enqueue(m_queue, nonce, g_thd)) {
-        return false;
-    }
-
-    if (!m_cn1->enqueue(m_queue, nonce, g_thd, w_size)) {
-        return false;
-    }
-
-    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;
-    }
+    enqueueReadBuffer(m_output, CL_TRUE, 0, sizeof(cl_uint) * 0x100, hashOutput);
 
     uint32_t &results = hashOutput[0xFF];
     if (results > 0xFF) {
@@ -168,38 +116,16 @@ bool xmrig::OclCnRunner::run(uint32_t nonce, uint32_t *hashOutput)
 }
 
 
-bool xmrig::OclCnRunner::selfTest() const
-{
-    if (OclBaseRunner::selfTest() && m_cn0->isValid() && m_cn2->isValid()) {
-        if (m_algorithm != Algorithm::CN_R) {
-            return m_cn1->isValid();
-        }
-
-        return true;
-    }
-
-    return false;
-}
-
-
 bool xmrig::OclCnRunner::set(const Job &job, uint8_t *blob)
 {
     if (job.size() > (Job::kMaxBlobSize - 4)) {
-        return false;
+        throw std::length_error("job size too big");
     }
 
     blob[job.size()] = 0x01;
     memset(blob + job.size() + 1, 0, Job::kMaxBlobSize - job.size() - 1);
 
-    if (OclLib::enqueueWriteBuffer(m_queue, m_input, CL_TRUE, 0, Job::kMaxBlobSize, blob, 0, nullptr, nullptr) != CL_SUCCESS) {
-        return false;
-    }
-
-    const uint32_t intensity = data().thread.intensity();
-
-    if (!m_cn0->setArgs(m_input, m_scratchpads, m_states, intensity)) {
-        return false;
-    }
+    enqueueWriteBuffer(m_input, CL_TRUE, 0, Job::kMaxBlobSize, blob);
 
     if (m_algorithm == Algorithm::CN_R && m_height != job.height()) {
         delete m_cn1;
@@ -207,20 +133,11 @@ bool xmrig::OclCnRunner::set(const Job &job, uint8_t *blob)
         m_height = job.height();
         m_cnr    = OclCnR::get(*this, m_height);
         m_cn1    = new Cn1Kernel(m_cnr, m_height);
+        m_cn1->setArgs(m_input, m_scratchpads, m_states, data().thread.intensity());
     }
 
-    if (!m_cn1->setArgs(m_input, m_scratchpads, m_states, intensity)) {
-        return false;
-    }
-
-    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;
-        }
+    for (auto kernel : m_branchKernels) {
+        kernel->setTarget(job.target());
     }
 
     return true;
@@ -231,18 +148,38 @@ void xmrig::OclCnRunner::build()
 {
     OclBaseRunner::build();
 
-    if (!m_program) {
-        return;
-    }
+    const uint32_t intensity = data().thread.intensity();
 
     m_cn0 = new Cn0Kernel(m_program);
+    m_cn0->setArgs(m_input, m_scratchpads, m_states, intensity);
+
     m_cn2 = new Cn2Kernel(m_program);
+    m_cn2->setArgs(m_scratchpads, m_states, m_branches, intensity);
 
     if (m_algorithm != Algorithm::CN_R) {
         m_cn1 = new Cn1Kernel(m_program);
+        m_cn1->setArgs(m_input, m_scratchpads, m_states, intensity);
     }
 
     for (size_t i = 0; i < BRANCH_MAX; ++i) {
-        m_branchKernels[i] = new CnBranchKernel(i, m_program);
+        auto kernel = new CnBranchKernel(i, m_program);
+        kernel->setArgs(m_states, m_branches[i], m_output, intensity);
+
+        m_branchKernels[i] = kernel;
+    }
+}
+
+
+void xmrig::OclCnRunner::init()
+{
+    OclBaseRunner::init();
+
+    const size_t g_thd = data().thread.intensity();
+
+    m_scratchpads = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, m_algorithm.l3() * g_thd);
+    m_states      = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, 200 * g_thd);
+
+    for (size_t i = 0; i < BRANCH_MAX; ++i) {
+        m_branches[i] = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, sizeof(cl_uint) * (g_thd + 2));
     }
 }
diff --git a/src/backend/opencl/runners/OclCnRunner.h b/src/backend/opencl/runners/OclCnRunner.h
index 34f6aad8..77c45b04 100644
--- a/src/backend/opencl/runners/OclCnRunner.h
+++ b/src/backend/opencl/runners/OclCnRunner.h
@@ -41,22 +41,16 @@ class CnBranchKernel;
 class OclCnRunner : public OclBaseRunner
 {
 public:
-    OclCnRunner()                         = delete;
-    OclCnRunner(const OclCnRunner &other) = delete;
-    OclCnRunner(OclCnRunner &&other)      = delete;
-    OclCnRunner(size_t index, const OclLaunchData &data);
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(OclCnRunner)
 
+    OclCnRunner(size_t index, const OclLaunchData &data);
     ~OclCnRunner() override;
 
-    OclCnRunner &operator=(const OclCnRunner &other) = delete;
-    OclCnRunner &operator=(OclCnRunner &&other)      = delete;
-
 protected:
-    bool isReadyToBuild() const override;
     bool run(uint32_t nonce, uint32_t *hashOutput) override;
-    bool selfTest() const override;
     bool set(const Job &job, uint8_t *blob) override;
     void build() override;
+    void init() override;
 
 private:
     enum Branches : size_t {
diff --git a/src/backend/opencl/runners/OclRxBaseRunner.cpp b/src/backend/opencl/runners/OclRxBaseRunner.cpp
new file mode 100644
index 00000000..84560e70
--- /dev/null
+++ b/src/backend/opencl/runners/OclRxBaseRunner.cpp
@@ -0,0 +1,129 @@
+/* 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/OclRxBaseRunner.h"
+
+#include "backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h"
+#include "backend/opencl/kernels/rx/Blake2bInitialHashKernel.h"
+#include "backend/opencl/kernels/rx/FillAesKernel.h"
+#include "backend/opencl/kernels/rx/FindSharesKernel.h"
+#include "backend/opencl/kernels/rx/HashAesKernel.h"
+#include "backend/opencl/OclLaunchData.h"
+#include "backend/opencl/wrappers/OclLib.h"
+#include "crypto/rx/RxAlgo.h"
+
+
+xmrig::OclRxBaseRunner::OclRxBaseRunner(size_t index, const OclLaunchData &data) : OclBaseRunner(index, data)
+{
+    uint32_t worksize    = 0;
+    uint32_t gcn_version = 12;
+
+    switch (data.thread.worksize()) {
+    case 2:
+    case 4:
+    case 8:
+    case 16:
+        worksize = data.thread.worksize();
+        break;
+
+    default:
+        worksize = 8;
+    }
+
+    if (data.device.type() == OclDevice::Vega_10 || data.device.type() == OclDevice::Vega_20) {
+        gcn_version = 14;
+    }
+
+    m_options += " -DALGO="             + std::to_string(m_algorithm.id());
+    m_options += " -DWORKERS_PER_HASH=" + std::to_string(worksize);
+    m_options += " -DGCN_VERSION="      + std::to_string(gcn_version);
+}
+
+
+xmrig::OclRxBaseRunner::~OclRxBaseRunner()
+{
+    delete m_fillAes1Rx4_scratchpad;
+    delete m_fillAes4Rx4_entropy;
+    delete m_hashAes1Rx4;
+    delete m_blake2b_initial_hash;
+    delete m_blake2b_hash_registers_32;
+    delete m_blake2b_hash_registers_64;
+    delete m_find_shares;
+
+    OclLib::release(m_entropy);
+    OclLib::release(m_hashes);
+    OclLib::release(m_rounding);
+    OclLib::release(m_scratchpads);
+}
+
+
+bool xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
+{
+    return false;
+}
+
+
+bool xmrig::OclRxBaseRunner::set(const Job &job, uint8_t *blob)
+{
+    return false;
+}
+
+
+void xmrig::OclRxBaseRunner::build()
+{
+    OclBaseRunner::build();
+
+    const uint32_t batch_size = data().thread.intensity();
+    const uint32_t rx_version = RxAlgo::version(m_algorithm);
+
+    m_fillAes1Rx4_scratchpad = new FillAesKernel(m_program, "fillAes1Rx4_scratchpad");
+    m_fillAes1Rx4_scratchpad->setArgs(m_hashes, m_scratchpads, batch_size, rx_version);
+
+    m_fillAes4Rx4_entropy = new FillAesKernel(m_program, "fillAes4Rx4_entropy");
+    m_fillAes1Rx4_scratchpad->setArgs(m_hashes, m_entropy, batch_size, rx_version);
+
+    m_hashAes1Rx4 = new HashAesKernel(m_program);
+
+    m_blake2b_initial_hash = new Blake2bInitialHashKernel(m_program);
+    m_blake2b_initial_hash->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");
+
+    m_find_shares = new FindSharesKernel(m_program);
+    m_find_shares->setArgs(m_hashes, m_output);
+}
+
+
+void xmrig::OclRxBaseRunner::init()
+{
+    OclBaseRunner::init();
+
+    const size_t g_thd = data().thread.intensity();
+
+    m_scratchpads = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, (m_algorithm.l3() + 64) * g_thd);
+    m_hashes      = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, 64 * g_thd);
+    m_entropy     = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, (128 + 2560) * g_thd);
+    m_rounding    = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, sizeof(uint32_t) * g_thd);
+}
diff --git a/src/backend/opencl/runners/OclRxBaseRunner.h b/src/backend/opencl/runners/OclRxBaseRunner.h
new file mode 100644
index 00000000..545bf8df
--- /dev/null
+++ b/src/backend/opencl/runners/OclRxBaseRunner.h
@@ -0,0 +1,74 @@
+/* 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_OCLRXBASERUNNER_H
+#define XMRIG_OCLRXBASERUNNER_H
+
+
+#include "backend/opencl/runners/OclBaseRunner.h"
+
+
+namespace xmrig {
+
+
+class Blake2bHashRegistersKernel;
+class Blake2bInitialHashKernel;
+class FillAesKernel;
+class FindSharesKernel;
+class HashAesKernel;
+
+
+class OclRxBaseRunner : public OclBaseRunner
+{
+public:
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(OclRxBaseRunner)
+
+    OclRxBaseRunner(size_t index, const OclLaunchData &data);
+    ~OclRxBaseRunner() override;
+
+protected:
+    bool run(uint32_t nonce, uint32_t *hashOutput) override;
+    bool set(const Job &job, uint8_t *blob) override;
+    void build() override;
+    void init() override;
+
+protected:
+    Blake2bHashRegistersKernel *m_blake2b_hash_registers_32 = nullptr;
+    Blake2bHashRegistersKernel *m_blake2b_hash_registers_64 = nullptr;
+    Blake2bInitialHashKernel *m_blake2b_initial_hash        = nullptr;
+    cl_mem m_entropy                                        = nullptr;
+    cl_mem m_hashes                                         = nullptr;
+    cl_mem m_rounding                                       = nullptr;
+    cl_mem m_scratchpads                                    = nullptr;
+    FillAesKernel *m_fillAes1Rx4_scratchpad                 = nullptr;
+    FillAesKernel *m_fillAes4Rx4_entropy                    = nullptr;
+    FindSharesKernel *m_find_shares                         = nullptr;
+    HashAesKernel *m_hashAes1Rx4                            = nullptr;
+};
+
+
+} /* namespace xmrig */
+
+
+#endif // XMRIG_OCLRXBASERUNNER_H
diff --git a/src/backend/opencl/runners/OclRxJitRunner.cpp b/src/backend/opencl/runners/OclRxJitRunner.cpp
new file mode 100644
index 00000000..00884d4f
--- /dev/null
+++ b/src/backend/opencl/runners/OclRxJitRunner.cpp
@@ -0,0 +1,76 @@
+/* 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/OclRxJitRunner.h"
+
+#include "backend/opencl/wrappers/OclLib.h"
+#include "backend/opencl/OclLaunchData.h"
+#include "backend/opencl/kernels/rx/HashAesKernel.h"
+#include "backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h"
+
+
+xmrig::OclRxJitRunner::OclRxJitRunner(size_t index, const OclLaunchData &data) : OclRxBaseRunner(index, data)
+{
+    if (m_rounding == nullptr) {
+        return;
+    }
+
+    const size_t g_thd = data.thread.intensity();
+    cl_int ret;
+
+    m_registers = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, 256 * g_thd, nullptr, &ret);
+    if (ret != CL_SUCCESS) {
+        return;
+    }
+
+    m_intermediate_programs = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, 5120 * g_thd, nullptr, &ret);
+    if (ret != CL_SUCCESS) {
+        return;
+    }
+
+    m_programs = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, 10048 * g_thd, nullptr, &ret);
+    if (ret != CL_SUCCESS) {
+        return;
+    }
+}
+
+
+xmrig::OclRxJitRunner::~OclRxJitRunner()
+{
+    OclLib::release(m_intermediate_programs);
+    OclLib::release(m_programs);
+    OclLib::release(m_registers);
+}
+
+
+void xmrig::OclRxJitRunner::build()
+{
+    OclRxBaseRunner::build();
+
+    const uint32_t batch_size = data().thread.intensity();
+
+    m_hashAes1Rx4->setArgs(m_scratchpads, m_registers, 256, batch_size);
+    m_blake2b_hash_registers_32->setArgs(m_hashes, m_registers, 256);
+    m_blake2b_hash_registers_64->setArgs(m_hashes, m_registers, 256);
+}
diff --git a/src/backend/opencl/runners/OclRxRunner.h b/src/backend/opencl/runners/OclRxJitRunner.h
similarity index 75%
rename from src/backend/opencl/runners/OclRxRunner.h
rename to src/backend/opencl/runners/OclRxJitRunner.h
index 345f16ae..6dc13173 100644
--- a/src/backend/opencl/runners/OclRxRunner.h
+++ b/src/backend/opencl/runners/OclRxJitRunner.h
@@ -22,26 +22,31 @@
  *   along with this program. If not, see <http://www.gnu.org/licenses/>.
  */
 
-#ifndef XMRIG_OCLRXRUNNER_H
-#define XMRIG_OCLRXRUNNER_H
+#ifndef XMRIG_OCLRXJITRUNNER_H
+#define XMRIG_OCLRXJITRUNNER_H
 
 
-#include "backend/opencl/runners/OclBaseRunner.h"
+#include "backend/opencl/runners/OclRxBaseRunner.h"
 
 
 namespace xmrig {
 
 
-class OclRxRunner : public OclBaseRunner
+class OclRxJitRunner : public OclRxBaseRunner
 {
 public:
-    OclRxRunner(size_t index, const OclLaunchData &data);
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(OclRxJitRunner)
+
+    OclRxJitRunner(size_t index, const OclLaunchData &data);
+    ~OclRxJitRunner() override;
 
 protected:
-    bool run(uint32_t nonce, uint32_t *hashOutput) override;
-    bool selfTest() const override;
-    bool set(const Job &job, uint8_t *blob) override;
     void build() override;
+
+private:
+    cl_mem m_intermediate_programs  = nullptr;
+    cl_mem m_programs               = nullptr;
+    cl_mem m_registers              = nullptr;
 };
 
 
diff --git a/src/backend/opencl/runners/OclRxVmRunner.cpp b/src/backend/opencl/runners/OclRxVmRunner.cpp
new file mode 100644
index 00000000..a15cac9a
--- /dev/null
+++ b/src/backend/opencl/runners/OclRxVmRunner.cpp
@@ -0,0 +1,77 @@
+/* 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/OclRxVmRunner.h"
+
+#include "backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h"
+#include "backend/opencl/kernels/rx/ExecuteVmKernel.h"
+#include "backend/opencl/kernels/rx/HashAesKernel.h"
+#include "backend/opencl/kernels/rx/InitVmKernel.h"
+#include "backend/opencl/OclLaunchData.h"
+#include "backend/opencl/wrappers/OclLib.h"
+#include "crypto/rx/RxAlgo.h"
+
+
+xmrig::OclRxVmRunner::OclRxVmRunner(size_t index, const OclLaunchData &data) : OclRxBaseRunner(index, data)
+{
+    if (m_rounding == nullptr) {
+        return;
+    }
+
+    const size_t g_thd = data.thread.intensity();
+    cl_int ret;
+
+    m_vm_states = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, 2560 * g_thd, nullptr, &ret);
+    if (ret != CL_SUCCESS) {
+        return;
+    }
+}
+
+
+xmrig::OclRxVmRunner::~OclRxVmRunner()
+{
+    delete m_init_vm;
+    delete m_execute_vm;
+
+    OclLib::release(m_vm_states);
+}
+
+
+void xmrig::OclRxVmRunner::build()
+{
+    OclRxBaseRunner::build();
+
+    const uint32_t batch_size       = data().thread.intensity();
+    const uint32_t hashStrideBytes  = RxAlgo::programSize(m_algorithm) * 8;
+
+    m_hashAes1Rx4->setArgs(m_scratchpads, m_vm_states, hashStrideBytes, batch_size);
+    m_blake2b_hash_registers_32->setArgs(m_hashes, m_vm_states, hashStrideBytes);
+    m_blake2b_hash_registers_64->setArgs(m_hashes, m_vm_states, hashStrideBytes);
+
+    m_init_vm = new InitVmKernel(m_program);
+    m_init_vm->setArgs(m_entropy, m_vm_states, m_rounding);
+
+    m_execute_vm = new ExecuteVmKernel(m_program);
+    m_execute_vm->setArgs(m_vm_states, m_rounding, m_scratchpads, data().dataset->get(), batch_size);
+}
diff --git a/src/backend/opencl/runners/OclRxVmRunner.h b/src/backend/opencl/runners/OclRxVmRunner.h
new file mode 100644
index 00000000..2d63583d
--- /dev/null
+++ b/src/backend/opencl/runners/OclRxVmRunner.h
@@ -0,0 +1,60 @@
+/* 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_OCLRXVMRUNNER_H
+#define XMRIG_OCLRXVMRUNNER_H
+
+
+#include "backend/opencl/runners/OclRxBaseRunner.h"
+
+
+namespace xmrig {
+
+
+class ExecuteVmKernel;
+class InitVmKernel;
+
+
+class OclRxVmRunner : public OclRxBaseRunner
+{
+public:
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(OclRxVmRunner)
+
+    OclRxVmRunner(size_t index, const OclLaunchData &data);
+    ~OclRxVmRunner() override;
+
+protected:
+    void build() override;
+
+private:
+    cl_mem m_vm_states              = nullptr;
+    ExecuteVmKernel *m_execute_vm   = nullptr;
+    InitVmKernel *m_init_vm         = nullptr;
+};
+
+
+} /* namespace xmrig */
+
+
+#endif // XMRIG_OCLRXVMRUNNER_H
diff --git a/src/backend/opencl/runners/OclRyoRunner.cpp b/src/backend/opencl/runners/OclRyoRunner.cpp
index 3d3183a3..9ae26b75 100644
--- a/src/backend/opencl/runners/OclRyoRunner.cpp
+++ b/src/backend/opencl/runners/OclRyoRunner.cpp
@@ -39,23 +39,6 @@
 
 xmrig::OclRyoRunner::OclRyoRunner(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);
-    if (ret != CL_SUCCESS) {
-        return;
-    }
-
     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";
@@ -78,14 +61,6 @@ xmrig::OclRyoRunner::~OclRyoRunner()
 }
 
 
-bool xmrig::OclRyoRunner::isReadyToBuild() const
-{
-    return OclBaseRunner::isReadyToBuild() &&
-            m_scratchpads                   != nullptr &&
-            m_states                        != nullptr;
-}
-
-
 bool xmrig::OclRyoRunner::run(uint32_t nonce, uint32_t *hashOutput)
 {
     static const cl_uint zero = 0;
@@ -96,29 +71,14 @@ bool xmrig::OclRyoRunner::run(uint32_t nonce, uint32_t *hashOutput)
 
     assert(g_thd % w_size == 0);
 
-    if (OclLib::enqueueWriteBuffer(m_queue, m_output, CL_FALSE, sizeof(cl_uint) * 0xFF, sizeof(cl_uint), &zero, 0, nullptr, nullptr) != CL_SUCCESS) {
-        return false;
-    }
+    enqueueWriteBuffer(m_output, CL_FALSE, sizeof(cl_uint) * 0xFF, sizeof(cl_uint), &zero);
 
-    if (!m_cn0->enqueue(m_queue, nonce, g_thd)) {
-        return false;
-    }
+    m_cn0->enqueue(m_queue, nonce, g_thd);
+    m_cn00->enqueue(m_queue, g_thd);
+    m_cn1->enqueue(m_queue, g_thd, w_size);
+    m_cn2->enqueue(m_queue, nonce, g_thd);
 
-    if (!m_cn00->enqueue(m_queue, g_thd)) {
-        return false;
-    }
-
-    if (!m_cn1->enqueue(m_queue, g_thd, w_size)) {
-        return false;
-    }
-
-    if (!m_cn2->enqueue(m_queue, nonce, g_thd)) {
-        return false;
-    }
-
-    if (OclLib::enqueueReadBuffer(m_queue, m_output, CL_TRUE, 0, sizeof(cl_uint) * 0x100, hashOutput, 0, nullptr, nullptr) != CL_SUCCESS) {
-        return false;
-    }
+    enqueueReadBuffer(m_output, CL_TRUE, 0, sizeof(cl_uint) * 0x100, hashOutput);
 
     uint32_t &results = hashOutput[0xFF];
     if (results > 0xFF) {
@@ -129,12 +89,6 @@ bool xmrig::OclRyoRunner::run(uint32_t nonce, uint32_t *hashOutput)
 }
 
 
-bool xmrig::OclRyoRunner::selfTest() const
-{
-    return OclBaseRunner::selfTest() && m_cn0->isValid() && m_cn00->isValid() && m_cn1->isValid() && m_cn2->isValid();
-}
-
-
 bool xmrig::OclRyoRunner::set(const Job &job, uint8_t *blob)
 {
     if (job.size() > (Job::kMaxBlobSize - 4)) {
@@ -144,27 +98,9 @@ bool xmrig::OclRyoRunner::set(const Job &job, uint8_t *blob)
     blob[job.size()] = 0x01;
     memset(blob + job.size() + 1, 0, Job::kMaxBlobSize - job.size() - 1);
 
-    if (OclLib::enqueueWriteBuffer(m_queue, m_input, CL_TRUE, 0, Job::kMaxBlobSize, blob, 0, nullptr, nullptr) != CL_SUCCESS) {
-        return false;
-    }
+    enqueueWriteBuffer(m_input, CL_TRUE, 0, Job::kMaxBlobSize, blob);
 
-    const uint32_t intensity = data().thread.intensity();
-
-    if (!m_cn00->setArgs(m_scratchpads, m_states)) {
-        return false;
-    }
-
-    if (!m_cn0->setArgs(m_input, m_scratchpads, m_states, intensity)) {
-        return false;
-    }
-
-    if (!m_cn1->setArgs(m_scratchpads, m_states, intensity)) {
-        return false;
-    }
-
-    if (!m_cn2->setArgs(m_scratchpads, m_states, m_output, job.target(), intensity)) {
-        return false;
-    }
+    m_cn2->setTarget(job.target());
 
     return true;
 }
@@ -174,12 +110,28 @@ void xmrig::OclRyoRunner::build()
 {
     OclBaseRunner::build();
 
-    if (!m_program) {
-        return;
-    }
+    const uint32_t intensity = data().thread.intensity();
 
     m_cn00 = new Cn00RyoKernel(m_program);
-    m_cn0  = new Cn0Kernel(m_program);
-    m_cn1  = new Cn1RyoKernel(m_program);
-    m_cn2  = new Cn2RyoKernel(m_program);
+    m_cn00->setArgs(m_scratchpads, m_states);
+
+    m_cn0 = new Cn0Kernel(m_program);
+    m_cn0->setArgs(m_input, m_scratchpads, m_states, intensity);
+
+    m_cn1 = new Cn1RyoKernel(m_program);
+    m_cn1->setArgs(m_scratchpads, m_states, intensity);
+
+    m_cn2 = new Cn2RyoKernel(m_program);
+    m_cn2->setArgs(m_scratchpads, m_states, m_output, intensity);
+}
+
+
+void xmrig::OclRyoRunner::init()
+{
+    OclBaseRunner::init();
+
+    const size_t g_thd = data().thread.intensity();
+
+    m_scratchpads = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, data().algorithm.l3() * g_thd);
+    m_states      = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, 200 * g_thd);
 }
diff --git a/src/backend/opencl/runners/OclRyoRunner.h b/src/backend/opencl/runners/OclRyoRunner.h
index ba4a93e3..f3f3c5e1 100644
--- a/src/backend/opencl/runners/OclRyoRunner.h
+++ b/src/backend/opencl/runners/OclRyoRunner.h
@@ -41,22 +41,17 @@ class Cn2RyoKernel;
 class OclRyoRunner : public OclBaseRunner
 {
 public:
-    OclRyoRunner()                          = delete;
-    OclRyoRunner(const OclRyoRunner &other) = delete;
-    OclRyoRunner(OclRyoRunner &&other)      = delete;
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(OclRyoRunner)
+
     OclRyoRunner(size_t index, const OclLaunchData &data);
 
     ~OclRyoRunner() override;
 
-    OclRyoRunner &operator=(const OclRyoRunner &other) = delete;
-    OclRyoRunner &operator=(OclRyoRunner &&other)      = delete;
-
 protected:
-    bool isReadyToBuild() const override;
     bool run(uint32_t nonce, uint32_t *hashOutput) override;
-    bool selfTest() const override;
     bool set(const Job &job, uint8_t *blob) override;
     void build() override;
+    void init() override;
 
 private:
     cl_mem m_scratchpads    = nullptr;
diff --git a/src/backend/opencl/runners/OclRxRunner.cpp b/src/backend/opencl/runners/tools/OclRxDataset.cpp
similarity index 52%
rename from src/backend/opencl/runners/OclRxRunner.cpp
rename to src/backend/opencl/runners/tools/OclRxDataset.cpp
index eb408df4..a3b54952 100644
--- a/src/backend/opencl/runners/OclRxRunner.cpp
+++ b/src/backend/opencl/runners/tools/OclRxDataset.cpp
@@ -22,61 +22,20 @@
  *   along with this program. If not, see <http://www.gnu.org/licenses/>.
  */
 
-#include "backend/opencl/runners/OclRxRunner.h"
 
-#include "backend/opencl/OclLaunchData.h"
+#include "backend/opencl/wrappers/OclLib.h"
+#include "backend/opencl/runners/tools/OclRxDataset.h"
+#include "crypto/rx/RxDataset.h"
 
 
-xmrig::OclRxRunner::OclRxRunner(size_t index, const OclLaunchData &data) : OclBaseRunner(index, data)
+void xmrig::OclRxDataset::createBuffer(cl_context ctx, const Algorithm &algorithm, bool host)
 {
-    uint32_t worksize    = 0;
-    uint32_t gcn_version = 12;
+    cl_int ret;
 
-    switch (data.thread.worksize()) {
-    case 2:
-    case 4:
-    case 8:
-    case 16:
-        worksize = data.thread.worksize();
-        break;
-
-    default:
-        worksize = 8;
+    if (host) {
+        // TODO use host memory for dataset
     }
-
-    if (data.device.type() == OclDevice::Vega_10 || data.device.type() == OclDevice::Vega_20) {
-        gcn_version = 14;
-    }
-
-    m_options += " -DALGO="             + std::to_string(m_algorithm.id());
-    m_options += " -DWORKERS_PER_HASH=" + std::to_string(worksize);
-    m_options += " -DGCN_VERSION="      + std::to_string(gcn_version);
-}
-
-
-bool xmrig::OclRxRunner::run(uint32_t nonce, uint32_t *hashOutput)
-{
-    return false;
-}
-
-
-bool xmrig::OclRxRunner::selfTest() const
-{
-    return false; // TODO OclRxRunner
-}
-
-
-bool xmrig::OclRxRunner::set(const Job &job, uint8_t *blob)
-{
-    return false;
-}
-
-
-void xmrig::OclRxRunner::build()
-{
-    OclBaseRunner::build();
-
-    if (!m_program) {
-        return;
+    else {
+        m_dataset = OclLib::createBuffer(ctx, CL_MEM_READ_ONLY, RxDataset::maxSize(), nullptr, &ret);
     }
 }
diff --git a/src/backend/opencl/runners/tools/OclRxDataset.h b/src/backend/opencl/runners/tools/OclRxDataset.h
new file mode 100644
index 00000000..198d4075
--- /dev/null
+++ b/src/backend/opencl/runners/tools/OclRxDataset.h
@@ -0,0 +1,62 @@
+/* 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_OCLRXDATASET_H
+#define XMRIG_OCLRXDATASET_H
+
+
+#include <memory>
+
+
+using cl_context = struct _cl_context *;
+using cl_mem     = struct _cl_mem *;
+
+
+namespace xmrig {
+
+
+class Algorithm;
+
+
+class OclRxDataset
+{
+public:
+    OclRxDataset() = default;
+
+    inline cl_mem get() const { return m_dataset; }
+
+    void createBuffer(cl_context ctx, const Algorithm &algorithm, bool host);
+
+private:
+    cl_mem m_dataset = nullptr;
+};
+
+
+using OclRxDatasetPtr = std::shared_ptr<OclRxDataset>;
+
+
+} /* namespace xmrig */
+
+
+#endif /* XMRIG_OCLINTERLEAVE_H */
diff --git a/src/backend/opencl/wrappers/OclContext.cpp b/src/backend/opencl/wrappers/OclContext.cpp
index bc0f3de4..e390478f 100644
--- a/src/backend/opencl/wrappers/OclContext.cpp
+++ b/src/backend/opencl/wrappers/OclContext.cpp
@@ -23,8 +23,8 @@
  */
 
 
-#include "backend/opencl/wrappers/OclLib.h"
 #include "backend/opencl/wrappers/OclContext.h"
+#include "backend/opencl/wrappers/OclLib.h"
 
 
 xmrig::OclContext::OclContext(const OclDevice &device)
@@ -59,6 +59,12 @@ bool xmrig::OclContext::init(const std::vector<OclDevice> &devices, std::vector<
 
     for (OclLaunchData &data : threads) {
         data.ctx = m_ctx;
+
+#       ifdef XMRIG_ALGO_RANDOMX
+        if (data.algorithm.family() == Algorithm::RANDOM_X) {
+            data.dataset->createBuffer(m_ctx, data.algorithm, data.thread.isDatasetHost());
+        }
+#       endif
     }
 
     return true;
diff --git a/src/backend/opencl/wrappers/OclContext.h b/src/backend/opencl/wrappers/OclContext.h
index d3107dac..f5f5380a 100644
--- a/src/backend/opencl/wrappers/OclContext.h
+++ b/src/backend/opencl/wrappers/OclContext.h
@@ -28,9 +28,10 @@
 
 #include "backend/opencl/OclLaunchData.h"
 #include "backend/opencl/wrappers/OclDevice.h"
+#include "base/tools/Object.h"
 
 
-typedef struct _cl_context *cl_context;
+using cl_context = struct _cl_context *;
 
 
 namespace xmrig {
@@ -39,6 +40,8 @@ namespace xmrig {
 class OclContext
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE(OclContext)
+
     OclContext() = default;
     OclContext(const OclDevice &device);
     ~OclContext();
diff --git a/src/backend/opencl/wrappers/OclKernel.cpp b/src/backend/opencl/wrappers/OclKernel.cpp
index 2da63458..7de93c47 100644
--- a/src/backend/opencl/wrappers/OclKernel.cpp
+++ b/src/backend/opencl/wrappers/OclKernel.cpp
@@ -23,17 +23,20 @@
  */
 
 
+#include "backend/common/Tags.h"
 #include "backend/opencl/wrappers/OclError.h"
 #include "backend/opencl/wrappers/OclKernel.h"
 #include "backend/opencl/wrappers/OclLib.h"
 #include "base/io/log/Log.h"
 
 
+#include <stdexcept>
+
+
 xmrig::OclKernel::OclKernel(cl_program program, const char *name) :
     m_name(name)
 {
-    cl_int ret = 0;
-    m_kernel = OclLib::createKernel(program, name, &ret);
+    m_kernel = OclLib::createKernel(program, name);
 }
 
 
@@ -43,34 +46,26 @@ xmrig::OclKernel::~OclKernel()
 }
 
 
-bool xmrig::OclKernel::enqueueNDRange(cl_command_queue queue, uint32_t work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size)
+void xmrig::OclKernel::enqueueNDRange(cl_command_queue queue, uint32_t work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size)
 {
-    if (!isValid()) {
-        return false;
-    }
-
     const cl_int ret = OclLib::enqueueNDRangeKernel(queue, m_kernel, work_dim, global_work_offset, global_work_size, local_work_size, 0, nullptr, nullptr);
     if (ret != CL_SUCCESS) {
-        LOG_ERR(MAGENTA_BG_BOLD(WHITE_BOLD_S " ocl ") RED(" error ") RED_BOLD("%s") RED(" when calling ") RED_BOLD("clEnqueueNDRangeKernel") RED(" for kernel ") RED_BOLD("%s"),
-                OclError::toString(ret), name().data());
-    }
+        LOG_ERR("%s" RED(" error ") RED_BOLD("%s") RED(" when calling ") RED_BOLD("clEnqueueNDRangeKernel") RED(" for kernel ") RED_BOLD("%s"),
+                ocl_tag(), OclError::toString(ret), name().data());
 
-    return ret == CL_SUCCESS;
+        throw std::runtime_error(OclError::toString(ret));
+    }
 }
 
 
-bool xmrig::OclKernel::setArg(uint32_t index, size_t size, const void *value)
+void 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")
+        LOG_ERR("%s" 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);
-    }
+                ocl_tag(), OclError::toString(ret), name().data(), index, size);
 
-    return ret == CL_SUCCESS;
+        throw std::runtime_error(OclError::toString(ret));
+    }
 }
diff --git a/src/backend/opencl/wrappers/OclKernel.h b/src/backend/opencl/wrappers/OclKernel.h
index d285c687..8664d130 100644
--- a/src/backend/opencl/wrappers/OclKernel.h
+++ b/src/backend/opencl/wrappers/OclKernel.h
@@ -26,6 +26,7 @@
 #define XMRIG_OCLKERNEL_H
 
 
+#include "base/tools/Object.h"
 #include "base/tools/String.h"
 
 
@@ -41,6 +42,8 @@ namespace xmrig {
 class OclKernel
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(OclKernel)
+
     OclKernel(cl_program program, const char *name);
     virtual ~OclKernel();
 
@@ -48,8 +51,8 @@ public:
     inline cl_kernel kernel() const     { return m_kernel; }
     inline const String &name() const   { return m_name; }
 
-    bool enqueueNDRange(cl_command_queue queue, uint32_t work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size);
-    bool setArg(uint32_t index, size_t size, const void *value);
+    void enqueueNDRange(cl_command_queue queue, uint32_t work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size);
+    void setArg(uint32_t index, size_t size, const void *value);
 
 private:
     cl_kernel m_kernel = nullptr;
diff --git a/src/backend/opencl/wrappers/OclLib.cpp b/src/backend/opencl/wrappers/OclLib.cpp
index 0c9a4b38..7ebe5a9c 100644
--- a/src/backend/opencl/wrappers/OclLib.cpp
+++ b/src/backend/opencl/wrappers/OclLib.cpp
@@ -27,6 +27,7 @@
 #include <uv.h>
 
 
+#include "backend/common/Tags.h"
 #include "backend/opencl/wrappers/OclError.h"
 #include "backend/opencl/wrappers/OclLib.h"
 #include "base/io/log/Log.h"
@@ -63,6 +64,7 @@ static const char *kReleaseMemObject                 = "clReleaseMemObject";
 static const char *kReleaseProgram                   = "clReleaseProgram";
 static const char *kSetKernelArg                     = "clSetKernelArg";
 
+
 #if defined(CL_VERSION_2_0)
 typedef cl_command_queue (CL_API_CALL *createCommandQueueWithProperties_t)(cl_context, cl_device_id, const cl_queue_properties *, cl_int *);
 #endif
@@ -208,7 +210,7 @@ const char *xmrig::OclLib::defaultLoader()
 }
 
 
-cl_command_queue xmrig::OclLib::createCommandQueue(cl_context context, cl_device_id device, cl_int *errcode_ret)
+cl_command_queue xmrig::OclLib::createCommandQueue(cl_context context, cl_device_id device, cl_int *errcode_ret) noexcept
 {
     cl_command_queue result;
 
@@ -235,6 +237,18 @@ cl_command_queue xmrig::OclLib::createCommandQueue(cl_context context, cl_device
 }
 
 
+cl_command_queue xmrig::OclLib::createCommandQueue(cl_context context, cl_device_id device)
+{
+    cl_int ret;
+    cl_command_queue queue = createCommandQueue(context, device, &ret);
+    if (ret != CL_SUCCESS) {
+        throw std::runtime_error(OclError::toString(ret));
+    }
+
+    return queue;
+}
+
+
 cl_context xmrig::OclLib::createContext(const cl_context_properties *properties, cl_uint num_devices, const cl_device_id *devices, void (CL_CALLBACK *pfn_notify)(const char *, const void *, size_t, void *), void *user_data, cl_int *errcode_ret)
 {
     assert(pCreateContext != nullptr);
@@ -257,7 +271,7 @@ cl_context xmrig::OclLib::createContext(const std::vector<cl_device_id> &ids)
 }
 
 
-cl_int xmrig::OclLib::buildProgram(cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void (CL_CALLBACK *pfn_notify)(cl_program program, void *user_data), void *user_data)
+cl_int xmrig::OclLib::buildProgram(cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void (CL_CALLBACK *pfn_notify)(cl_program program, void *user_data), void *user_data) noexcept
 {
     assert(pBuildProgram != nullptr);
 
@@ -270,7 +284,7 @@ cl_int xmrig::OclLib::buildProgram(cl_program program, cl_uint num_devices, cons
 }
 
 
-cl_int xmrig::OclLib::enqueueNDRangeKernel(cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
+cl_int xmrig::OclLib::enqueueNDRangeKernel(cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event) noexcept
 {
     assert(pEnqueueNDRangeKernel != nullptr);
 
@@ -278,7 +292,7 @@ cl_int xmrig::OclLib::enqueueNDRangeKernel(cl_command_queue command_queue, cl_ke
 }
 
 
-cl_int xmrig::OclLib::enqueueReadBuffer(cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_read, size_t offset, size_t size, void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
+cl_int xmrig::OclLib::enqueueReadBuffer(cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_read, size_t offset, size_t size, void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event) noexcept
 {
     assert(pEnqueueReadBuffer != nullptr);
 
@@ -291,7 +305,7 @@ cl_int xmrig::OclLib::enqueueReadBuffer(cl_command_queue command_queue, cl_mem b
 }
 
 
-cl_int xmrig::OclLib::enqueueWriteBuffer(cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write, size_t offset, size_t size, const void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
+cl_int xmrig::OclLib::enqueueWriteBuffer(cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write, size_t offset, size_t size, const void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event) noexcept
 {
     assert(pEnqueueWriteBuffer != nullptr);
 
@@ -304,7 +318,7 @@ cl_int xmrig::OclLib::enqueueWriteBuffer(cl_command_queue command_queue, cl_mem
 }
 
 
-cl_int xmrig::OclLib::finish(cl_command_queue command_queue)
+cl_int xmrig::OclLib::finish(cl_command_queue command_queue) noexcept
 {
     assert(pFinish != nullptr);
 
@@ -312,7 +326,7 @@ cl_int xmrig::OclLib::finish(cl_command_queue command_queue)
 }
 
 
-cl_int xmrig::OclLib::getDeviceIDs(cl_platform_id platform, cl_device_type device_type, cl_uint num_entries, cl_device_id *devices, cl_uint *num_devices)
+cl_int xmrig::OclLib::getDeviceIDs(cl_platform_id platform, cl_device_type device_type, cl_uint num_entries, cl_device_id *devices, cl_uint *num_devices) noexcept
 {
     assert(pGetDeviceIDs != nullptr);
 
@@ -320,7 +334,7 @@ cl_int xmrig::OclLib::getDeviceIDs(cl_platform_id platform, cl_device_type devic
 }
 
 
-cl_int xmrig::OclLib::getDeviceInfo(cl_device_id device, cl_device_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
+cl_int xmrig::OclLib::getDeviceInfo(cl_device_id device, cl_device_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret) noexcept
 {
     assert(pGetDeviceInfo != nullptr);
 
@@ -341,7 +355,7 @@ cl_int xmrig::OclLib::getPlatformIDs(cl_uint num_entries, cl_platform_id *platfo
 }
 
 
-cl_int xmrig::OclLib::getPlatformInfo(cl_platform_id platform, cl_platform_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
+cl_int xmrig::OclLib::getPlatformInfo(cl_platform_id platform, cl_platform_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret) noexcept
 {
     assert(pGetPlatformInfo != nullptr);
 
@@ -349,7 +363,7 @@ cl_int xmrig::OclLib::getPlatformInfo(cl_platform_id platform, cl_platform_info
 }
 
 
-cl_int xmrig::OclLib::getProgramBuildInfo(cl_program program, cl_device_id device, cl_program_build_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
+cl_int xmrig::OclLib::getProgramBuildInfo(cl_program program, cl_device_id device, cl_program_build_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret) noexcept
 {
     assert(pGetProgramBuildInfo != nullptr);
 
@@ -375,11 +389,15 @@ cl_int xmrig::OclLib::getProgramInfo(cl_program program, cl_program_info param_n
 }
 
 
-cl_int xmrig::OclLib::release(cl_command_queue command_queue)
+cl_int xmrig::OclLib::release(cl_command_queue command_queue) noexcept
 {
     assert(pReleaseCommandQueue != nullptr);
     assert(pGetCommandQueueInfo != nullptr);
 
+    if (command_queue == nullptr) {
+        return CL_SUCCESS;
+    }
+
     finish(command_queue);
 
     cl_int ret = pReleaseCommandQueue(command_queue);
@@ -391,7 +409,7 @@ cl_int xmrig::OclLib::release(cl_command_queue command_queue)
 }
 
 
-cl_int xmrig::OclLib::release(cl_context context)
+cl_int xmrig::OclLib::release(cl_context context) noexcept
 {
     assert(pReleaseContext != nullptr);
 
@@ -404,7 +422,7 @@ cl_int xmrig::OclLib::release(cl_context context)
 }
 
 
-cl_int xmrig::OclLib::release(cl_kernel kernel)
+cl_int xmrig::OclLib::release(cl_kernel kernel) noexcept
 {
     assert(pReleaseKernel != nullptr);
 
@@ -421,7 +439,7 @@ cl_int xmrig::OclLib::release(cl_kernel kernel)
 }
 
 
-cl_int xmrig::OclLib::release(cl_mem mem_obj)
+cl_int xmrig::OclLib::release(cl_mem mem_obj) noexcept
 {
     assert(pReleaseMemObject != nullptr);
 
@@ -438,7 +456,7 @@ cl_int xmrig::OclLib::release(cl_mem mem_obj)
 }
 
 
-cl_int xmrig::OclLib::release(cl_program program)
+cl_int xmrig::OclLib::release(cl_program program) noexcept
 {
     assert(pReleaseProgram != nullptr);
 
@@ -455,7 +473,7 @@ cl_int xmrig::OclLib::release(cl_program program)
 }
 
 
-cl_int xmrig::OclLib::setKernelArg(cl_kernel kernel, cl_uint arg_index, size_t arg_size, const void *arg_value)
+cl_int xmrig::OclLib::setKernelArg(cl_kernel kernel, cl_uint arg_index, size_t arg_size, const void *arg_value) noexcept
 {
     assert(pSetKernelArg != nullptr);
 
@@ -463,14 +481,14 @@ cl_int xmrig::OclLib::setKernelArg(cl_kernel kernel, cl_uint arg_index, size_t a
 }
 
 
-cl_kernel xmrig::OclLib::createKernel(cl_program program, const char *kernel_name, cl_int *errcode_ret)
+cl_kernel xmrig::OclLib::createKernel(cl_program program, const char *kernel_name, cl_int *errcode_ret) noexcept
 {
     assert(pCreateKernel != nullptr);
 
     auto result = pCreateKernel(program, kernel_name, 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("clCreateKernel") RED(" for kernel ") RED_BOLD("%s"),
-                OclError::toString(*errcode_ret), kernel_name);
+        LOG_ERR("%s" RED(" error ") RED_BOLD("%s") RED(" when calling ") RED_BOLD("clCreateKernel") RED(" for kernel ") RED_BOLD("%s"),
+                ocl_tag(), OclError::toString(*errcode_ret), kernel_name);
 
         return nullptr;
     }
@@ -479,14 +497,38 @@ cl_kernel xmrig::OclLib::createKernel(cl_program program, const char *kernel_nam
 }
 
 
-cl_mem xmrig::OclLib::createBuffer(cl_context context, cl_mem_flags flags, size_t size, void *host_ptr, cl_int *errcode_ret)
+cl_kernel xmrig::OclLib::createKernel(cl_program program, const char *kernel_name)
+{
+    cl_int ret;
+    cl_kernel kernel = createKernel(program, kernel_name, &ret);
+    if (ret != CL_SUCCESS) {
+        throw std::runtime_error(OclError::toString(ret));
+    }
+
+    return kernel;
+}
+
+
+cl_mem xmrig::OclLib::createBuffer(cl_context context, cl_mem_flags flags, size_t size, void *host_ptr)
+{
+    cl_int ret;
+    cl_mem mem = createBuffer(context, flags, size, host_ptr, &ret);
+    if (ret != CL_SUCCESS) {
+        throw std::runtime_error(OclError::toString(ret));
+    }
+
+    return mem;
+}
+
+
+cl_mem xmrig::OclLib::createBuffer(cl_context context, cl_mem_flags flags, size_t size, void *host_ptr, cl_int *errcode_ret) noexcept
 {
     assert(pCreateBuffer != nullptr);
 
     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);
+        LOG_ERR("%s" RED(" error ") RED_BOLD("%s") RED(" when calling ") RED_BOLD("%s") RED(" with buffer size ") RED_BOLD("%zu"),
+                ocl_tag(), OclError::toString(*errcode_ret), kCreateBuffer, size);
 
         return nullptr;
     }
@@ -495,7 +537,7 @@ cl_mem xmrig::OclLib::createBuffer(cl_context context, cl_mem_flags flags, size_
 }
 
 
-cl_program xmrig::OclLib::createProgramWithBinary(cl_context context, cl_uint num_devices, const cl_device_id *device_list, const size_t *lengths, const unsigned char **binaries, cl_int *binary_status, cl_int *errcode_ret)
+cl_program xmrig::OclLib::createProgramWithBinary(cl_context context, cl_uint num_devices, const cl_device_id *device_list, const size_t *lengths, const unsigned char **binaries, cl_int *binary_status, cl_int *errcode_ret) noexcept
 {
     assert(pCreateProgramWithBinary != nullptr);
 
@@ -510,7 +552,7 @@ cl_program xmrig::OclLib::createProgramWithBinary(cl_context context, cl_uint nu
 }
 
 
-cl_program xmrig::OclLib::createProgramWithSource(cl_context context, cl_uint count, const char **strings, const size_t *lengths, cl_int *errcode_ret)
+cl_program xmrig::OclLib::createProgramWithSource(cl_context context, cl_uint count, const char **strings, const size_t *lengths, cl_int *errcode_ret) noexcept
 {
     assert(pCreateProgramWithSource != nullptr);
 
@@ -525,7 +567,7 @@ cl_program xmrig::OclLib::createProgramWithSource(cl_context context, cl_uint co
 }
 
 
-cl_uint xmrig::OclLib::getDeviceUint(cl_device_id id, cl_device_info param, cl_uint defaultValue)
+cl_uint xmrig::OclLib::getDeviceUint(cl_device_id id, cl_device_info param, cl_uint defaultValue) noexcept
 {
     OclLib::getDeviceInfo(id, param, sizeof(cl_uint), &defaultValue);
 
@@ -533,7 +575,7 @@ cl_uint xmrig::OclLib::getDeviceUint(cl_device_id id, cl_device_info param, cl_u
 }
 
 
-cl_uint xmrig::OclLib::getNumPlatforms()
+cl_uint xmrig::OclLib::getNumPlatforms() noexcept
 {
     cl_uint count = 0;
     cl_int ret;
@@ -550,7 +592,7 @@ cl_uint xmrig::OclLib::getNumPlatforms()
 }
 
 
-cl_uint xmrig::OclLib::getReferenceCount(cl_program program)
+cl_uint xmrig::OclLib::getReferenceCount(cl_program program) noexcept
 {
     cl_uint out = 0;
     OclLib::getProgramInfo(program, CL_PROGRAM_REFERENCE_COUNT, sizeof(cl_uint), &out);
@@ -559,7 +601,7 @@ cl_uint xmrig::OclLib::getReferenceCount(cl_program program)
 }
 
 
-cl_ulong xmrig::OclLib::getDeviceUlong(cl_device_id id, cl_device_info param, cl_ulong defaultValue)
+cl_ulong xmrig::OclLib::getDeviceUlong(cl_device_id id, cl_device_info param, cl_ulong defaultValue) noexcept
 {
     OclLib::getDeviceInfo(id, param, sizeof(cl_ulong), &defaultValue);
 
@@ -567,7 +609,7 @@ cl_ulong xmrig::OclLib::getDeviceUlong(cl_device_id id, cl_device_info param, cl
 }
 
 
-std::vector<cl_platform_id> xmrig::OclLib::getPlatformIDs()
+std::vector<cl_platform_id> xmrig::OclLib::getPlatformIDs() noexcept
 {
     const uint32_t count = getNumPlatforms();
     std::vector<cl_platform_id> platforms(count);
@@ -580,7 +622,7 @@ std::vector<cl_platform_id> xmrig::OclLib::getPlatformIDs()
 }
 
 
-xmrig::String xmrig::OclLib::getDeviceString(cl_device_id id, cl_device_info param)
+xmrig::String xmrig::OclLib::getDeviceString(cl_device_id id, cl_device_info param) noexcept
 {
     size_t size = 0;
     if (getDeviceInfo(id, param, 0, nullptr, &size) != CL_SUCCESS) {
@@ -594,7 +636,7 @@ xmrig::String xmrig::OclLib::getDeviceString(cl_device_id id, cl_device_info par
 }
 
 
-xmrig::String xmrig::OclLib::getPlatformInfo(cl_platform_id platform, cl_platform_info param_name)
+xmrig::String xmrig::OclLib::getPlatformInfo(cl_platform_id platform, cl_platform_info param_name) noexcept
 {
     size_t size = 0;
     if (getPlatformInfo(platform, param_name, 0, nullptr, &size) != CL_SUCCESS) {
@@ -608,7 +650,7 @@ xmrig::String xmrig::OclLib::getPlatformInfo(cl_platform_id platform, cl_platfor
 }
 
 
-xmrig::String xmrig::OclLib::getProgramBuildLog(cl_program program, cl_device_id device)
+xmrig::String xmrig::OclLib::getProgramBuildLog(cl_program program, cl_device_id device) noexcept
 {
     size_t size = 0;
     if (getProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, nullptr, &size) != CL_SUCCESS) {
diff --git a/src/backend/opencl/wrappers/OclLib.h b/src/backend/opencl/wrappers/OclLib.h
index aa85a4e6..887514f6 100644
--- a/src/backend/opencl/wrappers/OclLib.h
+++ b/src/backend/opencl/wrappers/OclLib.h
@@ -46,38 +46,41 @@ public:
     static inline bool isInitialized()   { return m_initialized; }
     static inline const String &loader() { return m_loader; }
 
-    static cl_command_queue createCommandQueue(cl_context context, cl_device_id device, cl_int *errcode_ret);
+    static cl_command_queue createCommandQueue(cl_context context, cl_device_id device, cl_int *errcode_ret) noexcept;
+    static cl_command_queue createCommandQueue(cl_context context, cl_device_id device);
     static cl_context createContext(const cl_context_properties *properties, cl_uint num_devices, const cl_device_id *devices, void (CL_CALLBACK *pfn_notify)(const char *, const void *, size_t, void *), void *user_data, cl_int *errcode_ret);
     static cl_context createContext(const std::vector<cl_device_id> &ids);
-    static cl_int buildProgram(cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options = nullptr, void (CL_CALLBACK *pfn_notify)(cl_program program, void *user_data) = nullptr, void *user_data = nullptr);
-    static cl_int enqueueNDRangeKernel(cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event);
-    static cl_int enqueueReadBuffer(cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_read, size_t offset, size_t size, void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event);
-    static cl_int enqueueWriteBuffer(cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write, size_t offset, size_t size, const void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event);
-    static cl_int finish(cl_command_queue command_queue);
-    static cl_int getDeviceIDs(cl_platform_id platform, cl_device_type device_type, cl_uint num_entries, cl_device_id *devices, cl_uint *num_devices);
-    static cl_int getDeviceInfo(cl_device_id device, cl_device_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret = nullptr);
+    static cl_int buildProgram(cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options = nullptr, void (CL_CALLBACK *pfn_notify)(cl_program program, void *user_data) = nullptr, void *user_data = nullptr) noexcept;
+    static cl_int enqueueNDRangeKernel(cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event) noexcept;
+    static cl_int enqueueReadBuffer(cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_read, size_t offset, size_t size, void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event) noexcept;
+    static cl_int enqueueWriteBuffer(cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write, size_t offset, size_t size, const void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event) noexcept;
+    static cl_int finish(cl_command_queue command_queue) noexcept;
+    static cl_int getDeviceIDs(cl_platform_id platform, cl_device_type device_type, cl_uint num_entries, cl_device_id *devices, cl_uint *num_devices) noexcept;
+    static cl_int getDeviceInfo(cl_device_id device, cl_device_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret = nullptr) noexcept;
     static cl_int getPlatformIDs(cl_uint num_entries, cl_platform_id *platforms, cl_uint *num_platforms);
-    static cl_int getPlatformInfo(cl_platform_id platform, cl_platform_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret);
-    static cl_int getProgramBuildInfo(cl_program program, cl_device_id device, cl_program_build_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret);
+    static cl_int getPlatformInfo(cl_platform_id platform, cl_platform_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret) noexcept;
+    static cl_int getProgramBuildInfo(cl_program program, cl_device_id device, cl_program_build_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret) noexcept;
     static cl_int getProgramInfo(cl_program program, cl_program_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret = nullptr);
-    static cl_int release(cl_command_queue command_queue);
-    static cl_int release(cl_context context);
-    static cl_int release(cl_kernel kernel);
-    static cl_int release(cl_mem mem_obj);
-    static cl_int release(cl_program program);
-    static cl_int setKernelArg(cl_kernel kernel, cl_uint arg_index, size_t arg_size, const void *arg_value);
-    static cl_kernel createKernel(cl_program program, const char *kernel_name, cl_int *errcode_ret);
-    static cl_mem createBuffer(cl_context context, cl_mem_flags flags, size_t size, void *host_ptr, cl_int *errcode_ret);
-    static cl_program createProgramWithBinary(cl_context context, cl_uint num_devices, const cl_device_id *device_list, const size_t *lengths, const unsigned char **binaries, cl_int *binary_status, cl_int *errcode_ret);
-    static cl_program createProgramWithSource(cl_context context, cl_uint count, const char **strings, const size_t *lengths, cl_int *errcode_ret);
-    static cl_uint getDeviceUint(cl_device_id id, cl_device_info param, cl_uint defaultValue = 0);
-    static cl_uint getNumPlatforms();
-    static cl_uint getReferenceCount(cl_program program);
-    static cl_ulong getDeviceUlong(cl_device_id id, cl_device_info param, cl_ulong defaultValue = 0);
-    static std::vector<cl_platform_id> getPlatformIDs();
-    static String getDeviceString(cl_device_id id, cl_device_info param);
-    static String getPlatformInfo(cl_platform_id platform, cl_platform_info param_name);
-    static String getProgramBuildLog(cl_program program, cl_device_id device);
+    static cl_int release(cl_command_queue command_queue) noexcept;
+    static cl_int release(cl_context context) noexcept;
+    static cl_int release(cl_kernel kernel) noexcept;
+    static cl_int release(cl_mem mem_obj) noexcept;
+    static cl_int release(cl_program program) noexcept;
+    static cl_int setKernelArg(cl_kernel kernel, cl_uint arg_index, size_t arg_size, const void *arg_value) noexcept;
+    static cl_kernel createKernel(cl_program program, const char *kernel_name, cl_int *errcode_ret) noexcept;
+    static cl_kernel createKernel(cl_program program, const char *kernel_name);
+    static cl_mem createBuffer(cl_context context, cl_mem_flags flags, size_t size, void *host_ptr = nullptr);
+    static cl_mem createBuffer(cl_context context, cl_mem_flags flags, size_t size, void *host_ptr, cl_int *errcode_ret) noexcept;
+    static cl_program createProgramWithBinary(cl_context context, cl_uint num_devices, const cl_device_id *device_list, const size_t *lengths, const unsigned char **binaries, cl_int *binary_status, cl_int *errcode_ret) noexcept;
+    static cl_program createProgramWithSource(cl_context context, cl_uint count, const char **strings, const size_t *lengths, cl_int *errcode_ret) noexcept;
+    static cl_uint getDeviceUint(cl_device_id id, cl_device_info param, cl_uint defaultValue = 0) noexcept;
+    static cl_uint getNumPlatforms() noexcept;
+    static cl_uint getReferenceCount(cl_program program) noexcept;
+    static cl_ulong getDeviceUlong(cl_device_id id, cl_device_info param, cl_ulong defaultValue = 0) noexcept;
+    static std::vector<cl_platform_id> getPlatformIDs() noexcept;
+    static String getDeviceString(cl_device_id id, cl_device_info param) noexcept;
+    static String getPlatformInfo(cl_platform_id platform, cl_platform_info param_name) noexcept;
+    static String getProgramBuildLog(cl_program program, cl_device_id device) noexcept;
 
 private:
     static bool load();
diff --git a/src/base/tools/Object.h b/src/base/tools/Object.h
new file mode 100644
index 00000000..7e460e44
--- /dev/null
+++ b/src/base/tools/Object.h
@@ -0,0 +1,52 @@
+/* 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_OBJECT_H
+#define XMRIG_OBJECT_H
+
+
+#include <chrono>
+
+
+namespace xmrig {
+
+
+#define XMRIG_DISABLE_COPY_MOVE(X) \
+    X(const X &other)            = delete; \
+    X(X &&other)                 = delete; \
+    X &operator=(const X &other) = delete; \
+    X &operator=(X &&other)      = delete;
+
+
+#define XMRIG_DISABLE_COPY_MOVE_DEFAULT(X) \
+    X()                          = delete; \
+    X(const X &other)            = delete; \
+    X(X &&other)                 = delete; \
+    X &operator=(const X &other) = delete; \
+    X &operator=(X &&other)      = delete;
+
+
+} /* namespace xmrig */
+
+#endif /* XMRIG_OBJECT_H */
diff --git a/src/crypto/rx/Rx.cpp b/src/crypto/rx/Rx.cpp
index 2f9cb4d2..81c39f67 100644
--- a/src/crypto/rx/Rx.cpp
+++ b/src/crypto/rx/Rx.cpp
@@ -25,17 +25,7 @@
  */
 
 
-#include <map>
-#include <mutex>
-#include <thread>
-#include <uv.h>
-
-
-#ifdef XMRIG_FEATURE_HWLOC
-#   include <hwloc.h>
-#   include "backend/cpu/platform/HwlocCpuInfo.h"
-#endif
-
+#include "crypto/rx/Rx.h"
 
 #include "backend/common/interfaces/IRxListener.h"
 #include "backend/cpu/Cpu.h"
@@ -45,12 +35,24 @@
 #include "base/tools/Buffer.h"
 #include "base/tools/Chrono.h"
 #include "base/tools/Handle.h"
-#include "crypto/rx/Rx.h"
+#include "base/tools/Object.h"
 #include "crypto/rx/RxAlgo.h"
 #include "crypto/rx/RxCache.h"
 #include "crypto/rx/RxDataset.h"
 
 
+#ifdef XMRIG_FEATURE_HWLOC
+#   include <hwloc.h>
+#   include "backend/cpu/platform/HwlocCpuInfo.h"
+#endif
+
+
+#include <map>
+#include <mutex>
+#include <thread>
+#include <uv.h>
+
+
 namespace xmrig {
 
 
@@ -92,8 +94,9 @@ inline static void bindToNUMANode(uint32_t) {}
 class RxPrivate
 {
 public:
-    inline RxPrivate() :
-        m_seed()
+    XMRIG_DISABLE_COPY_MOVE(RxPrivate)
+
+    inline RxPrivate()
     {
         m_async = new uv_async_t;
         m_async->data = this;
@@ -144,12 +147,12 @@ public:
         LOG_INFO("%s" CYAN_BOLD("#%u") MAGENTA_BOLD(" allocate") CYAN_BOLD(" %zu MB") BLACK_BOLD(" (%zu+%zu) for RandomX dataset & cache"),
                  tag,
                  nodeId,
-                 (RxDataset::size() + RxCache::size()) / 1024 / 1024,
-                 RxDataset::size() / 1024 / 1024,
-                 RxCache::size() / 1024 / 1024
+                 (RxDataset::maxSize() + RxCache::maxSize()) / 1024 / 1024,
+                 RxDataset::maxSize() / 1024 / 1024,
+                 RxCache::maxSize() / 1024 / 1024
                  );
 
-        RxDataset *dataset      = new RxDataset(d_ptr->m_hugePages);
+        auto dataset            = new RxDataset(d_ptr->m_hugePages);
         d_ptr->datasets[nodeId] = dataset;
 
         if (dataset->get() != nullptr) {
@@ -244,7 +247,7 @@ private:
     bool m_numa             = true;
     IRxListener *m_listener = nullptr;
     size_t m_ready          = 0;
-    uint8_t m_seed[32];
+    uint8_t m_seed[32]{ 0 };
     uv_async_t *m_async;
 };
 
diff --git a/src/crypto/rx/Rx.h b/src/crypto/rx/Rx.h
index cdff9504..9392e80d 100644
--- a/src/crypto/rx/Rx.h
+++ b/src/crypto/rx/Rx.h
@@ -28,7 +28,7 @@
 #define XMRIG_RX_H
 
 
-#include <stdint.h>
+#include <cstdint>
 #include <utility>
 
 
diff --git a/src/crypto/rx/RxAlgo.cpp b/src/crypto/rx/RxAlgo.cpp
index daf3ec2d..b306a4bc 100644
--- a/src/crypto/rx/RxAlgo.cpp
+++ b/src/crypto/rx/RxAlgo.cpp
@@ -47,3 +47,29 @@ xmrig::Algorithm::Id xmrig::RxAlgo::apply(Algorithm::Id algorithm)
 
     return algorithm;
 }
+
+
+uint32_t xmrig::RxAlgo::version(Algorithm::Id algorithm)
+{
+    return algorithm == Algorithm::RX_WOW ? 103 : 104;
+}
+
+
+uint32_t xmrig::RxAlgo::programSize(Algorithm::Id algorithm)
+{
+    switch (algorithm) {
+    case Algorithm::RX_0:
+        return RandomX_MoneroConfig.ProgramSize;
+
+    case Algorithm::RX_WOW:
+        return RandomX_WowneroConfig.ProgramSize;
+
+    case Algorithm::RX_LOKI:
+        return RandomX_LokiConfig.ProgramSize;
+
+    default:
+        break;
+    }
+
+    return 0;
+}
diff --git a/src/crypto/rx/RxAlgo.h b/src/crypto/rx/RxAlgo.h
index 95033a9c..859e488d 100644
--- a/src/crypto/rx/RxAlgo.h
+++ b/src/crypto/rx/RxAlgo.h
@@ -28,8 +28,8 @@
 #define XMRIG_RX_ALGO_H
 
 
-#include <stddef.h>
-#include <stdint.h>
+#include <cstddef>
+#include <cstdint>
 
 
 #include "crypto/common/Algorithm.h"
@@ -43,6 +43,8 @@ class RxAlgo
 {
 public:
     static Algorithm::Id apply(Algorithm::Id algorithm);
+    static uint32_t programSize(Algorithm::Id algorithm);
+    static uint32_t version(Algorithm::Id algorithm);
 };
 
 
diff --git a/src/crypto/rx/RxCache.h b/src/crypto/rx/RxCache.h
index e6b2397c..12d173a2 100644
--- a/src/crypto/rx/RxCache.h
+++ b/src/crypto/rx/RxCache.h
@@ -28,7 +28,7 @@
 #define XMRIG_RX_CACHE_H
 
 
-#include <stdint.h>
+#include <cstdint>
 
 
 #include "crypto/randomx/configuration.h"
@@ -55,7 +55,7 @@ public:
 
     bool init(const uint8_t *seed);
 
-    static inline constexpr size_t size() { return RANDOMX_CACHE_MAX_SIZE; }
+    static inline constexpr size_t maxSize() { return RANDOMX_CACHE_MAX_SIZE; }
 
 private:
     bool isReady(const uint8_t *seed) const;
diff --git a/src/crypto/rx/RxDataset.cpp b/src/crypto/rx/RxDataset.cpp
index 50459a55..3373abee 100644
--- a/src/crypto/rx/RxDataset.cpp
+++ b/src/crypto/rx/RxDataset.cpp
@@ -99,16 +99,16 @@ bool xmrig::RxDataset::init(const uint8_t *seed, uint32_t numThreads)
 std::pair<size_t, size_t> xmrig::RxDataset::hugePages() const
 {
     constexpr size_t twoMiB      = 2u * 1024u * 1024u;
-    constexpr const size_t total = (VirtualMemory::align(size(), twoMiB) + VirtualMemory::align(RxCache::size(), twoMiB)) / twoMiB;
+    constexpr const size_t total = (VirtualMemory::align(maxSize(), twoMiB) + VirtualMemory::align(RxCache::maxSize(), twoMiB)) / twoMiB;
 
     size_t count = 0;
     if (isHugePages()) {
-        count += VirtualMemory::align(size(), twoMiB) / twoMiB;
+        count += VirtualMemory::align(maxSize(), twoMiB) / twoMiB;
     }
 
     if (m_cache->isHugePages()) {
-        count += VirtualMemory::align(RxCache::size(), twoMiB) / twoMiB;
+        count += VirtualMemory::align(RxCache::maxSize(), twoMiB) / twoMiB;
     }
 
-    return std::pair<size_t, size_t>(count, total);
+    return { count, total };
 }
diff --git a/src/crypto/rx/RxDataset.h b/src/crypto/rx/RxDataset.h
index 932f4ed9..491827a3 100644
--- a/src/crypto/rx/RxDataset.h
+++ b/src/crypto/rx/RxDataset.h
@@ -30,6 +30,7 @@
 
 #include "crypto/common/Algorithm.h"
 #include "crypto/randomx/configuration.h"
+#include "base/tools/Object.h"
 
 
 struct randomx_dataset;
@@ -45,6 +46,8 @@ class RxCache;
 class RxDataset
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(RxDataset)
+
     RxDataset(bool hugePages = true);
     ~RxDataset();
 
@@ -55,7 +58,7 @@ public:
     bool init(const uint8_t *seed, uint32_t numThreads);
     std::pair<size_t, size_t> hugePages() const;
 
-    static inline constexpr size_t size() { return RANDOMX_DATASET_MAX_SIZE; }
+    static inline constexpr size_t maxSize() { return RANDOMX_DATASET_MAX_SIZE; }
 
 private:
     Algorithm m_algorithm;