diff --git a/src/backend/cpu/CpuWorker.cpp b/src/backend/cpu/CpuWorker.cpp
index ed871f4a..a98027c0 100644
--- a/src/backend/cpu/CpuWorker.cpp
+++ b/src/backend/cpu/CpuWorker.cpp
@@ -24,7 +24,7 @@
  */
 
 
-#include <assert.h>
+#include <cassert>
 #include <thread>
 
 
diff --git a/src/backend/cpu/CpuWorker.h b/src/backend/cpu/CpuWorker.h
index 1bfb46a6..02ce0826 100644
--- a/src/backend/cpu/CpuWorker.h
+++ b/src/backend/cpu/CpuWorker.h
@@ -30,6 +30,7 @@
 #include "backend/common/Worker.h"
 #include "backend/common/WorkerJob.h"
 #include "backend/cpu/CpuLaunchData.h"
+#include "base/tools/Object.h"
 #include "net/JobResult.h"
 
 
@@ -43,6 +44,8 @@ template<size_t N>
 class CpuWorker : public Worker
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(CpuWorker)
+
     CpuWorker(size_t id, const CpuLaunchData &data);
     ~CpuWorker() override;
 
@@ -70,7 +73,7 @@ private:
     const CnHash::AlgoVariant m_av;
     const Miner *m_miner;
     cryptonight_ctx *m_ctx[N];
-    uint8_t m_hash[N * 32];
+    uint8_t m_hash[N * 32]{ 0 };
     VirtualMemory *m_memory = nullptr;
     WorkerJob<N> m_job;
 
diff --git a/src/backend/opencl/interfaces/IOclRunner.h b/src/backend/opencl/interfaces/IOclRunner.h
index c432ba0a..005c1b08 100644
--- a/src/backend/opencl/interfaces/IOclRunner.h
+++ b/src/backend/opencl/interfaces/IOclRunner.h
@@ -51,8 +51,8 @@ public:
     IOclRunner()          = default;
     virtual ~IOclRunner() = default;
 
-    virtual bool run(uint32_t nonce, uint32_t *hashOutput)  = 0;
-    virtual bool set(const Job &job, uint8_t *blob)         = 0;
+    virtual void run(uint32_t nonce, uint32_t *hashOutput)  = 0;
+    virtual void set(const Job &job, uint8_t *blob)         = 0;
     virtual cl_context ctx() const                          = 0;
     virtual const Algorithm &algorithm() const              = 0;
     virtual const char *buildOptions() const                = 0;
diff --git a/src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.cpp b/src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.cpp
index aea009e5..1b60d31b 100644
--- a/src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.cpp
+++ b/src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.cpp
@@ -27,6 +27,15 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
+void xmrig::Blake2bHashRegistersKernel::enqueue(cl_command_queue queue, size_t threads)
+{
+    const size_t gthreads        = threads;
+    static const size_t lthreads = 64;
+
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+}
+
+
 // __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)
diff --git a/src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h b/src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h
index f1f316c4..538033a0 100644
--- a/src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h
+++ b/src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h
@@ -37,6 +37,7 @@ class Blake2bHashRegistersKernel : public OclKernel
 public:
     inline Blake2bHashRegistersKernel(cl_program program, const char *name) : OclKernel(program, name) {}
 
+    void enqueue(cl_command_queue queue, size_t threads);
     void setArgs(cl_mem out, cl_mem in, uint32_t inStrideBytes);
 };
 
diff --git a/src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.cpp b/src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.cpp
index 334da909..7b6126e0 100644
--- a/src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.cpp
+++ b/src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.cpp
@@ -27,9 +27,32 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
+void xmrig::Blake2bInitialHashKernel::enqueue(cl_command_queue queue, size_t threads)
+{
+    const size_t gthreads        = threads;
+    static const size_t lthreads = 64;
+
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+}
+
+
 // __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);
 }
+
+
+void xmrig::Blake2bInitialHashKernel::setBlobSize(size_t size)
+{
+    const uint32_t s = size;
+
+    setArg(2, sizeof(uint32_t), &s);
+}
+
+
+void xmrig::Blake2bInitialHashKernel::setNonce(uint32_t nonce)
+{
+    setArg(3, sizeof(uint32_t), &nonce);
+}
diff --git a/src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.h b/src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.h
index 0eba796b..2eb6627b 100644
--- a/src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.h
+++ b/src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.h
@@ -37,7 +37,10 @@ class Blake2bInitialHashKernel : public OclKernel
 public:
     inline Blake2bInitialHashKernel(cl_program program) : OclKernel(program, "blake2b_initial_hash") {}
 
+    void enqueue(cl_command_queue queue, size_t threads);
     void setArgs(cl_mem out, cl_mem blockTemplate);
+    void setBlobSize(size_t size);
+    void setNonce(uint32_t nonce);
 };
 
 
diff --git a/src/backend/opencl/kernels/rx/ExecuteVmKernel.cpp b/src/backend/opencl/kernels/rx/ExecuteVmKernel.cpp
index c319757b..0b2968b1 100644
--- a/src/backend/opencl/kernels/rx/ExecuteVmKernel.cpp
+++ b/src/backend/opencl/kernels/rx/ExecuteVmKernel.cpp
@@ -27,6 +27,15 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
+void xmrig::ExecuteVmKernel::enqueue(cl_command_queue queue, size_t threads, size_t worksize)
+{
+    const size_t gthreads = (worksize == 16) ? (threads * 16) : (threads * 8);
+    const size_t lthreads = (worksize == 16) ? 32 : 16;
+
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+}
+
+
 // __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)
 {
@@ -36,3 +45,21 @@ void xmrig::ExecuteVmKernel::setArgs(cl_mem vm_states, cl_mem rounding, cl_mem s
     setArg(3, sizeof(cl_mem), &dataset_ptr);
     setArg(4, sizeof(uint32_t), &batch_size);
 }
+
+
+void xmrig::ExecuteVmKernel::setFirst(uint32_t first)
+{
+    setArg(6, sizeof(uint32_t), &first);
+}
+
+
+void xmrig::ExecuteVmKernel::setIterations(uint32_t num_iterations)
+{
+    setArg(5, sizeof(uint32_t), &num_iterations);
+}
+
+
+void xmrig::ExecuteVmKernel::setLast(uint32_t last)
+{
+    setArg(7, sizeof(uint32_t), &last);
+}
diff --git a/src/backend/opencl/kernels/rx/ExecuteVmKernel.h b/src/backend/opencl/kernels/rx/ExecuteVmKernel.h
index 3f9215ef..4ab5813f 100644
--- a/src/backend/opencl/kernels/rx/ExecuteVmKernel.h
+++ b/src/backend/opencl/kernels/rx/ExecuteVmKernel.h
@@ -37,7 +37,11 @@ class ExecuteVmKernel : public OclKernel
 public:
     inline ExecuteVmKernel(cl_program program) : OclKernel(program, "execute_vm") {}
 
+    void enqueue(cl_command_queue queue, size_t threads, size_t worksize);
     void setArgs(cl_mem vm_states, cl_mem rounding, cl_mem scratchpads, cl_mem dataset_ptr, uint32_t batch_size);
+    void setFirst(uint32_t first);
+    void setIterations(uint32_t num_iterations);
+    void setLast(uint32_t last);
 };
 
 
diff --git a/src/backend/opencl/kernels/rx/FillAesKernel.cpp b/src/backend/opencl/kernels/rx/FillAesKernel.cpp
index ec960bfa..b196b406 100644
--- a/src/backend/opencl/kernels/rx/FillAesKernel.cpp
+++ b/src/backend/opencl/kernels/rx/FillAesKernel.cpp
@@ -27,6 +27,15 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
+void xmrig::FillAesKernel::enqueue(cl_command_queue queue, size_t threads)
+{
+    const size_t gthreads        = threads * 4;
+    static const size_t lthreads = 64;
+
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+}
+
+
 // __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)
diff --git a/src/backend/opencl/kernels/rx/FillAesKernel.h b/src/backend/opencl/kernels/rx/FillAesKernel.h
index 8d84a706..59d7e59c 100644
--- a/src/backend/opencl/kernels/rx/FillAesKernel.h
+++ b/src/backend/opencl/kernels/rx/FillAesKernel.h
@@ -37,6 +37,7 @@ class FillAesKernel : public OclKernel
 public:
     inline FillAesKernel(cl_program program, const char *name) : OclKernel(program, name) {}
 
+    void enqueue(cl_command_queue queue, size_t threads);
     void setArgs(cl_mem state, cl_mem out, uint32_t batch_size, uint32_t rx_version);
 };
 
diff --git a/src/backend/opencl/kernels/rx/FindSharesKernel.cpp b/src/backend/opencl/kernels/rx/FindSharesKernel.cpp
index fa50619e..583112ed 100644
--- a/src/backend/opencl/kernels/rx/FindSharesKernel.cpp
+++ b/src/backend/opencl/kernels/rx/FindSharesKernel.cpp
@@ -27,9 +27,30 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
+void xmrig::FindSharesKernel::enqueue(cl_command_queue queue, size_t threads)
+{
+    const size_t gthreads        = threads;
+    static const size_t lthreads = 64;
+
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+}
+
+
 // __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);
 }
+
+
+void xmrig::FindSharesKernel::setTarget(uint64_t target)
+{
+    setArg(1, sizeof(uint64_t), &target);
+}
+
+
+void xmrig::FindSharesKernel::setNonce(uint32_t nonce)
+{
+    setArg(2, sizeof(uint32_t), &nonce);
+}
diff --git a/src/backend/opencl/kernels/rx/FindSharesKernel.h b/src/backend/opencl/kernels/rx/FindSharesKernel.h
index 3c10284e..b1fe1427 100644
--- a/src/backend/opencl/kernels/rx/FindSharesKernel.h
+++ b/src/backend/opencl/kernels/rx/FindSharesKernel.h
@@ -37,7 +37,10 @@ class FindSharesKernel : public OclKernel
 public:
     inline FindSharesKernel(cl_program program) : OclKernel(program, "find_shares") {}
 
+    void enqueue(cl_command_queue queue, size_t threads);
     void setArgs(cl_mem hashes, cl_mem shares);
+    void setTarget(uint64_t target);
+    void setNonce(uint32_t nonce);
 };
 
 
diff --git a/src/backend/opencl/kernels/rx/HashAesKernel.cpp b/src/backend/opencl/kernels/rx/HashAesKernel.cpp
index 3a512269..3bdbc8bc 100644
--- a/src/backend/opencl/kernels/rx/HashAesKernel.cpp
+++ b/src/backend/opencl/kernels/rx/HashAesKernel.cpp
@@ -27,6 +27,15 @@
 #include "backend/opencl/wrappers/OclLib.h"
 
 
+void xmrig::HashAesKernel::enqueue(cl_command_queue queue, size_t threads)
+{
+    const size_t gthreads        = threads * 4;
+    static const size_t lthreads = 64;
+
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+}
+
+
 // __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)
 {
diff --git a/src/backend/opencl/kernels/rx/HashAesKernel.h b/src/backend/opencl/kernels/rx/HashAesKernel.h
index 33d9f17d..41026239 100644
--- a/src/backend/opencl/kernels/rx/HashAesKernel.h
+++ b/src/backend/opencl/kernels/rx/HashAesKernel.h
@@ -37,6 +37,7 @@ class HashAesKernel : public OclKernel
 public:
     inline HashAesKernel(cl_program program) : OclKernel(program, "hashAes1Rx4") {}
 
+    void enqueue(cl_command_queue queue, size_t threads);
     void setArgs(cl_mem input, cl_mem hash, uint32_t hashStrideBytes, uint32_t batch_size);
 };
 
diff --git a/src/backend/opencl/kernels/rx/InitVmKernel.cpp b/src/backend/opencl/kernels/rx/InitVmKernel.cpp
index 2c083fd5..a897922b 100644
--- a/src/backend/opencl/kernels/rx/InitVmKernel.cpp
+++ b/src/backend/opencl/kernels/rx/InitVmKernel.cpp
@@ -26,6 +26,24 @@
 #include "backend/opencl/kernels/rx/InitVmKernel.h"
 #include "backend/opencl/wrappers/OclLib.h"
 
+#include "base/io/log/Log.h"
+#include <thread>
+
+
+void xmrig::InitVmKernel::enqueue(cl_command_queue queue, size_t threads, uint32_t iteration)
+{
+    setArg(3, sizeof(uint32_t), &iteration);
+
+    const size_t gthreads        = threads * 8;
+    static const size_t lthreads = 32;
+
+//    LOG_WARN("%zu %zu %u", gthreads, lthreads, iteration);
+
+//    std::this_thread::sleep_for(std::chrono::milliseconds(500));
+
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
+}
+
 
 // __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)
diff --git a/src/backend/opencl/kernels/rx/InitVmKernel.h b/src/backend/opencl/kernels/rx/InitVmKernel.h
index b1fd6f8a..bba7cf43 100644
--- a/src/backend/opencl/kernels/rx/InitVmKernel.h
+++ b/src/backend/opencl/kernels/rx/InitVmKernel.h
@@ -37,6 +37,7 @@ class InitVmKernel : public OclKernel
 public:
     inline InitVmKernel(cl_program program) : OclKernel(program, "init_vm") {}
 
+    void enqueue(cl_command_queue queue, size_t threads, uint32_t iteration);
     void setArgs(cl_mem entropy_data, cl_mem vm_states, cl_mem rounding);
 };
 
diff --git a/src/backend/opencl/runners/OclBaseRunner.cpp b/src/backend/opencl/runners/OclBaseRunner.cpp
index 94c6a68c..32b0e017 100644
--- a/src/backend/opencl/runners/OclBaseRunner.cpp
+++ b/src/backend/opencl/runners/OclBaseRunner.cpp
@@ -105,3 +105,14 @@ void xmrig::OclBaseRunner::enqueueWriteBuffer(cl_mem buffer, cl_bool blocking_wr
         throw std::runtime_error(OclError::toString(ret));
     }
 }
+
+
+void xmrig::OclBaseRunner::finalize(uint32_t *hashOutput)
+{
+    enqueueReadBuffer(m_output, CL_TRUE, 0, sizeof(cl_uint) * 0x100, hashOutput);
+
+    uint32_t &results = hashOutput[0xFF];
+    if (results > 0xFF) {
+        results = 0xFF;
+    }
+}
diff --git a/src/backend/opencl/runners/OclBaseRunner.h b/src/backend/opencl/runners/OclBaseRunner.h
index ee3d4b5b..530bdce9 100644
--- a/src/backend/opencl/runners/OclBaseRunner.h
+++ b/src/backend/opencl/runners/OclBaseRunner.h
@@ -64,6 +64,7 @@ protected:
 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);
+    void finalize(uint32_t *hashOutput);
 
     Algorithm m_algorithm;
     cl_command_queue m_queue    = nullptr;
diff --git a/src/backend/opencl/runners/OclCnRunner.cpp b/src/backend/opencl/runners/OclCnRunner.cpp
index 49c60099..ab5aff15 100644
--- a/src/backend/opencl/runners/OclCnRunner.cpp
+++ b/src/backend/opencl/runners/OclCnRunner.cpp
@@ -81,7 +81,7 @@ xmrig::OclCnRunner::~OclCnRunner()
 }
 
 
-bool xmrig::OclCnRunner::run(uint32_t nonce, uint32_t *hashOutput)
+void xmrig::OclCnRunner::run(uint32_t nonce, uint32_t *hashOutput)
 {
     static const cl_uint zero = 0;
 
@@ -105,18 +105,11 @@ bool xmrig::OclCnRunner::run(uint32_t nonce, uint32_t *hashOutput)
         kernel->enqueue(m_queue, nonce, g_thd, w_size);
     }
 
-    enqueueReadBuffer(m_output, CL_TRUE, 0, sizeof(cl_uint) * 0x100, hashOutput);
-
-    uint32_t &results = hashOutput[0xFF];
-    if (results > 0xFF) {
-        results = 0xFF;
-    }
-
-    return true;
+    finalize(hashOutput);
 }
 
 
-bool xmrig::OclCnRunner::set(const Job &job, uint8_t *blob)
+void xmrig::OclCnRunner::set(const Job &job, uint8_t *blob)
 {
     if (job.size() > (Job::kMaxBlobSize - 4)) {
         throw std::length_error("job size too big");
@@ -139,8 +132,6 @@ bool xmrig::OclCnRunner::set(const Job &job, uint8_t *blob)
     for (auto kernel : m_branchKernels) {
         kernel->setTarget(job.target());
     }
-
-    return true;
 }
 
 
diff --git a/src/backend/opencl/runners/OclCnRunner.h b/src/backend/opencl/runners/OclCnRunner.h
index 77c45b04..e08496c0 100644
--- a/src/backend/opencl/runners/OclCnRunner.h
+++ b/src/backend/opencl/runners/OclCnRunner.h
@@ -47,8 +47,8 @@ public:
     ~OclCnRunner() override;
 
 protected:
-    bool run(uint32_t nonce, uint32_t *hashOutput) override;
-    bool set(const Job &job, uint8_t *blob) override;
+    void run(uint32_t nonce, uint32_t *hashOutput) override;
+    void set(const Job &job, uint8_t *blob) override;
     void build() override;
     void init() override;
 
diff --git a/src/backend/opencl/runners/OclRxBaseRunner.cpp b/src/backend/opencl/runners/OclRxBaseRunner.cpp
index 84560e70..6eab7c2c 100644
--- a/src/backend/opencl/runners/OclRxBaseRunner.cpp
+++ b/src/backend/opencl/runners/OclRxBaseRunner.cpp
@@ -31,7 +31,10 @@
 #include "backend/opencl/kernels/rx/HashAesKernel.h"
 #include "backend/opencl/OclLaunchData.h"
 #include "backend/opencl/wrappers/OclLib.h"
+#include "base/net/stratum/Job.h"
+#include "crypto/rx/Rx.h"
 #include "crypto/rx/RxAlgo.h"
+#include "crypto/rx/RxDataset.h"
 
 
 xmrig::OclRxBaseRunner::OclRxBaseRunner(size_t index, const OclLaunchData &data) : OclBaseRunner(index, data)
@@ -78,15 +81,61 @@ xmrig::OclRxBaseRunner::~OclRxBaseRunner()
 }
 
 
-bool xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
+void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
 {
-    return false;
+    static const uint32_t zero = 0;
+
+    m_blake2b_initial_hash->setNonce(nonce);
+    m_find_shares->setNonce(nonce);
+
+    enqueueWriteBuffer(m_output, CL_FALSE, sizeof(cl_uint) * 0xFF, sizeof(uint32_t), &zero);
+
+    const uint32_t g_intensity = data().thread.intensity();
+
+    m_blake2b_initial_hash->enqueue(m_queue, g_intensity);
+    m_fillAes1Rx4_scratchpad->enqueue(m_queue, g_intensity);
+
+    const uint32_t programCount = RxAlgo::programCount(m_algorithm);
+
+    for (uint32_t i = 0; i < programCount; ++i) {
+        m_fillAes4Rx4_entropy->enqueue(m_queue, g_intensity);
+
+        execute(i);
+
+        if (i == programCount - 1) {
+            m_hashAes1Rx4->enqueue(m_queue, g_intensity);
+            m_blake2b_hash_registers_32->enqueue(m_queue, g_intensity);
+        }
+        else {
+            m_blake2b_hash_registers_64->enqueue(m_queue, g_intensity);
+        }
+    }
+
+    m_find_shares->enqueue(m_queue, g_intensity);
+
+    finalize(hashOutput);
+
+    OclLib::finish(m_queue);
 }
 
 
-bool xmrig::OclRxBaseRunner::set(const Job &job, uint8_t *blob)
+void xmrig::OclRxBaseRunner::set(const Job &job, uint8_t *blob)
 {
-    return false;
+    if (!data().thread.isDatasetHost() && m_seed != job.seed()) {
+        m_seed = job.seed();
+
+        auto dataset = Rx::dataset(job, 0);
+        enqueueWriteBuffer(data().dataset->get(), CL_TRUE, 0, dataset->size(), dataset->raw());
+    }
+
+    if (job.size() < Job::kMaxBlobSize) {
+        memset(blob + job.size(), 0, Job::kMaxBlobSize - job.size());
+    }
+
+    enqueueWriteBuffer(m_input, CL_TRUE, 0, Job::kMaxBlobSize, blob);
+
+    m_blake2b_initial_hash->setBlobSize(job.size());
+    m_find_shares->setTarget(job.target());
 }
 
 
@@ -101,7 +150,7 @@ void xmrig::OclRxBaseRunner::build()
     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_fillAes4Rx4_entropy->setArgs(m_hashes, m_entropy, batch_size, rx_version);
 
     m_hashAes1Rx4 = new HashAesKernel(m_program);
 
diff --git a/src/backend/opencl/runners/OclRxBaseRunner.h b/src/backend/opencl/runners/OclRxBaseRunner.h
index 545bf8df..996a3c07 100644
--- a/src/backend/opencl/runners/OclRxBaseRunner.h
+++ b/src/backend/opencl/runners/OclRxBaseRunner.h
@@ -27,6 +27,7 @@
 
 
 #include "backend/opencl/runners/OclBaseRunner.h"
+#include "base/tools/Buffer.h"
 
 
 namespace xmrig {
@@ -48,15 +49,18 @@ public:
     ~OclRxBaseRunner() override;
 
 protected:
-    bool run(uint32_t nonce, uint32_t *hashOutput) override;
-    bool set(const Job &job, uint8_t *blob) override;
+    void run(uint32_t nonce, uint32_t *hashOutput) override;
+    void set(const Job &job, uint8_t *blob) override;
     void build() override;
     void init() override;
 
 protected:
+    virtual void execute(uint32_t iteration) = 0;
+
     Blake2bHashRegistersKernel *m_blake2b_hash_registers_32 = nullptr;
     Blake2bHashRegistersKernel *m_blake2b_hash_registers_64 = nullptr;
     Blake2bInitialHashKernel *m_blake2b_initial_hash        = nullptr;
+    Buffer m_seed;
     cl_mem m_entropy                                        = nullptr;
     cl_mem m_hashes                                         = nullptr;
     cl_mem m_rounding                                       = nullptr;
diff --git a/src/backend/opencl/runners/OclRxJitRunner.cpp b/src/backend/opencl/runners/OclRxJitRunner.cpp
index 00884d4f..da51e398 100644
--- a/src/backend/opencl/runners/OclRxJitRunner.cpp
+++ b/src/backend/opencl/runners/OclRxJitRunner.cpp
@@ -32,27 +32,6 @@
 
 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;
-    }
 }
 
 
@@ -74,3 +53,20 @@ void xmrig::OclRxJitRunner::build()
     m_blake2b_hash_registers_32->setArgs(m_hashes, m_registers, 256);
     m_blake2b_hash_registers_64->setArgs(m_hashes, m_registers, 256);
 }
+
+
+void xmrig::OclRxJitRunner::execute(uint32_t iteration)
+{
+}
+
+
+void xmrig::OclRxJitRunner::init()
+{
+    OclRxBaseRunner::init();
+
+    const size_t g_thd = data().thread.intensity();
+
+    m_registers             = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, 256 * g_thd, nullptr);
+    m_intermediate_programs = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, 5120 * g_thd, nullptr);
+    m_programs              = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, 10048 * g_thd, nullptr);
+}
diff --git a/src/backend/opencl/runners/OclRxJitRunner.h b/src/backend/opencl/runners/OclRxJitRunner.h
index 6dc13173..1e110b01 100644
--- a/src/backend/opencl/runners/OclRxJitRunner.h
+++ b/src/backend/opencl/runners/OclRxJitRunner.h
@@ -42,6 +42,8 @@ public:
 
 protected:
     void build() override;
+    void execute(uint32_t iteration) override;
+    void init() override;
 
 private:
     cl_mem m_intermediate_programs  = nullptr;
diff --git a/src/backend/opencl/runners/OclRxVmRunner.cpp b/src/backend/opencl/runners/OclRxVmRunner.cpp
index a15cac9a..898fc20d 100644
--- a/src/backend/opencl/runners/OclRxVmRunner.cpp
+++ b/src/backend/opencl/runners/OclRxVmRunner.cpp
@@ -32,20 +32,11 @@
 #include "backend/opencl/wrappers/OclLib.h"
 #include "crypto/rx/RxAlgo.h"
 
+#include "base/io/log/Log.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;
-    }
 }
 
 
@@ -75,3 +66,44 @@ void xmrig::OclRxVmRunner::build()
     m_execute_vm = new ExecuteVmKernel(m_program);
     m_execute_vm->setArgs(m_vm_states, m_rounding, m_scratchpads, data().dataset->get(), batch_size);
 }
+
+
+void xmrig::OclRxVmRunner::execute(uint32_t iteration)
+{
+    const uint32_t bfactor        = std::min(data().thread.bfactor(), 8u);
+    const uint32_t num_iterations = RxAlgo::programIterations(m_algorithm) >> bfactor;
+    const uint32_t g_intensity    = data().thread.intensity();
+
+    m_init_vm->enqueue(m_queue, g_intensity, iteration);
+
+//    LOG_WARN("bfactor:%u %u %u", bfactor, RxAlgo::programIterations(m_algorithm), num_iterations);
+
+    uint32_t first = 1;
+    uint32_t last  = 0;
+
+    m_execute_vm->setIterations(num_iterations);
+    m_execute_vm->setFirst(first);
+    m_execute_vm->setLast(last);
+
+    for (int j = 0, n = 1 << bfactor; j < n; ++j) {
+        if (j == n - 1) {
+            last = 1;
+            m_execute_vm->setLast(last);
+        }
+
+        m_execute_vm->enqueue(m_queue, g_intensity, data().thread.worksize());
+
+        if (j == 0) {
+            first = 0;
+            m_execute_vm->setFirst(first);
+        }
+    }
+}
+
+
+void xmrig::OclRxVmRunner::init()
+{
+    OclRxBaseRunner::init();
+
+    m_vm_states = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, 2560 * data().thread.intensity());
+}
diff --git a/src/backend/opencl/runners/OclRxVmRunner.h b/src/backend/opencl/runners/OclRxVmRunner.h
index 2d63583d..a1cbd188 100644
--- a/src/backend/opencl/runners/OclRxVmRunner.h
+++ b/src/backend/opencl/runners/OclRxVmRunner.h
@@ -46,6 +46,8 @@ public:
 
 protected:
     void build() override;
+    void execute(uint32_t iteration) override;
+    void init() override;
 
 private:
     cl_mem m_vm_states              = nullptr;
diff --git a/src/backend/opencl/runners/OclRyoRunner.cpp b/src/backend/opencl/runners/OclRyoRunner.cpp
index 9ae26b75..1b3d08c7 100644
--- a/src/backend/opencl/runners/OclRyoRunner.cpp
+++ b/src/backend/opencl/runners/OclRyoRunner.cpp
@@ -61,7 +61,7 @@ xmrig::OclRyoRunner::~OclRyoRunner()
 }
 
 
-bool xmrig::OclRyoRunner::run(uint32_t nonce, uint32_t *hashOutput)
+void xmrig::OclRyoRunner::run(uint32_t nonce, uint32_t *hashOutput)
 {
     static const cl_uint zero = 0;
 
@@ -78,21 +78,14 @@ bool xmrig::OclRyoRunner::run(uint32_t nonce, uint32_t *hashOutput)
     m_cn1->enqueue(m_queue, g_thd, w_size);
     m_cn2->enqueue(m_queue, nonce, g_thd);
 
-    enqueueReadBuffer(m_output, CL_TRUE, 0, sizeof(cl_uint) * 0x100, hashOutput);
-
-    uint32_t &results = hashOutput[0xFF];
-    if (results > 0xFF) {
-        results = 0xFF;
-    }
-
-    return true;
+    finalize(hashOutput);
 }
 
 
-bool xmrig::OclRyoRunner::set(const Job &job, uint8_t *blob)
+void xmrig::OclRyoRunner::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;
@@ -101,8 +94,6 @@ bool xmrig::OclRyoRunner::set(const Job &job, uint8_t *blob)
     enqueueWriteBuffer(m_input, CL_TRUE, 0, Job::kMaxBlobSize, blob);
 
     m_cn2->setTarget(job.target());
-
-    return true;
 }
 
 
diff --git a/src/backend/opencl/runners/OclRyoRunner.h b/src/backend/opencl/runners/OclRyoRunner.h
index f3f3c5e1..e236a3d6 100644
--- a/src/backend/opencl/runners/OclRyoRunner.h
+++ b/src/backend/opencl/runners/OclRyoRunner.h
@@ -48,8 +48,8 @@ public:
     ~OclRyoRunner() override;
 
 protected:
-    bool run(uint32_t nonce, uint32_t *hashOutput) override;
-    bool set(const Job &job, uint8_t *blob) override;
+    void run(uint32_t nonce, uint32_t *hashOutput) override;
+    void set(const Job &job, uint8_t *blob) override;
     void build() override;
     void init() override;
 
diff --git a/src/base/net/stratum/Job.cpp b/src/base/net/stratum/Job.cpp
index 04bd82d8..c86ecba5 100644
--- a/src/base/net/stratum/Job.cpp
+++ b/src/base/net/stratum/Job.cpp
@@ -25,32 +25,18 @@
  */
 
 
-#include <assert.h>
-#include <string.h>
+#include <cassert>
+#include <cstring>
 
 
 #include "base/net/stratum/Job.h"
 #include "base/tools/Buffer.h"
 
 
-xmrig::Job::Job() :
-    m_blob(),
-    m_seedHash()
-{
-}
-
-
 xmrig::Job::Job(bool nicehash, const Algorithm &algorithm, const String &clientId) :
     m_algorithm(algorithm),
     m_nicehash(nicehash),
-    m_clientId(clientId),
-    m_blob(),
-    m_seedHash()
-{
-}
-
-
-xmrig::Job::~Job()
+    m_clientId(clientId)
 {
 }
 
@@ -96,7 +82,7 @@ bool xmrig::Job::setBlob(const char *blob)
 
 bool xmrig::Job::setSeedHash(const char *hash)
 {
-    if (!hash || (strlen(hash) != sizeof(m_seedHash) * 2)) {
+    if (!hash || (strlen(hash) != kMaxSeedSize * 2)) {
         return false;
     }
 
@@ -104,7 +90,9 @@ bool xmrig::Job::setSeedHash(const char *hash)
     m_rawSeedHash = hash;
 #   endif
 
-    return Buffer::fromHex(hash, sizeof(m_seedHash) * 2, m_seedHash);
+    m_seed = Buffer::fromHex(hash, kMaxSeedSize * 2);
+
+    return !m_seed.isEmpty();
 }
 
 
@@ -173,9 +161,9 @@ void xmrig::Job::copy(const Job &other)
     m_height    = other.m_height;
     m_target    = other.m_target;
     m_index     = other.m_index;
+    m_seed      = other.m_seed;
 
     memcpy(m_blob, other.m_blob, sizeof(m_blob));
-    memcpy(m_seedHash, other.m_seedHash, sizeof(m_seedHash));
 
 #   ifdef XMRIG_PROXY_PROJECT
     m_rawSeedHash = other.m_rawSeedHash;
diff --git a/src/base/net/stratum/Job.h b/src/base/net/stratum/Job.h
index 2b256a12..eafd10e3 100644
--- a/src/base/net/stratum/Job.h
+++ b/src/base/net/stratum/Job.h
@@ -28,10 +28,11 @@
 #define XMRIG_JOB_H
 
 
-#include <stddef.h>
-#include <stdint.h>
+#include <cstddef>
+#include <cstdint>
 
 
+#include "base/tools/Buffer.h"
 #include "base/tools/String.h"
 #include "crypto/common/Algorithm.h"
 
@@ -45,10 +46,11 @@ public:
     // Max blob size is 84 (75 fixed + 9 variable), aligned to 96. https://github.com/xmrig/xmrig/issues/1 Thanks fireice-uk.
     // SECOR increase requirements for blob size: https://github.com/xmrig/xmrig/issues/913
     static constexpr const size_t kMaxBlobSize = 128;
+    static constexpr const size_t kMaxSeedSize = 32;
 
-    Job();
+    Job() = default;
     Job(bool nicehash, const Algorithm &algorithm, const String &clientId);
-    ~Job();
+    ~Job() = default;
 
     bool isEqual(const Job &other) const;
     bool setBlob(const char *blob);
@@ -60,11 +62,11 @@ public:
     inline bool isValid() const                       { return m_size > 0 && m_diff > 0; }
     inline bool setId(const char *id)                 { return m_id = id; }
     inline const Algorithm &algorithm() const         { return m_algorithm; }
+    inline const Buffer &seed() const                 { return m_seed; }
     inline const String &clientId() const             { return m_clientId; }
     inline const String &id() const                   { return m_id; }
     inline const uint32_t *nonce() const              { return reinterpret_cast<const uint32_t*>(m_blob + 39); }
     inline const uint8_t *blob() const                { return m_blob; }
-    inline const uint8_t *seedHash() const            { return m_seedHash; }
     inline size_t size() const                        { return m_size; }
     inline uint32_t *nonce()                          { return reinterpret_cast<uint32_t*>(m_blob + 39); }
     inline uint64_t diff() const                      { return m_diff; }
@@ -97,15 +99,15 @@ private:
 
     Algorithm m_algorithm;
     bool m_nicehash     = false;
+    Buffer m_seed;
     size_t m_size       = 0;
     String m_clientId;
     String m_id;
     uint64_t m_diff     = 0;
     uint64_t m_height   = 0;
     uint64_t m_target   = 0;
-    uint8_t m_blob[kMaxBlobSize];
+    uint8_t m_blob[kMaxBlobSize]{ 0 };
     uint8_t m_index     = 0;
-    uint8_t m_seedHash[32];
 
 #   ifdef XMRIG_PROXY_PROJECT
     char m_rawBlob[kMaxBlobSize * 2 + 8];
diff --git a/src/base/tools/Buffer.cpp b/src/base/tools/Buffer.cpp
index 05a68cb4..c12e601d 100644
--- a/src/base/tools/Buffer.cpp
+++ b/src/base/tools/Buffer.cpp
@@ -53,14 +53,7 @@ static inline uint8_t hf_bin2hex(uint8_t c)
 }
 
 
-xmrig::Buffer::Buffer() :
-    m_data(nullptr),
-    m_size(0)
-{
-}
-
-
-xmrig::Buffer::Buffer(Buffer &&other) :
+xmrig::Buffer::Buffer(Buffer &&other) noexcept :
     m_data(other.m_data),
     m_size(other.m_size)
 {
@@ -138,11 +131,13 @@ bool xmrig::Buffer::fromHex(const uint8_t *in, size_t size, uint8_t *out)
 xmrig::Buffer xmrig::Buffer::fromHex(const char *data, size_t size)
 {
     if (data == nullptr || size % 2 != 0) {
-        return Buffer();
+        return {};
     }
 
     Buffer buf(size / 2);
-    fromHex(data, size, buf.data());
+    if (!fromHex(data, size, buf.data())) {
+        return {};
+    }
 
     return buf;
 }
@@ -157,12 +152,6 @@ void xmrig::Buffer::toHex(const uint8_t *in, size_t size, uint8_t *out)
 }
 
 
-xmrig::String xmrig::Buffer::toHex(const uint8_t *in, size_t size)
-{
-    return Buffer(reinterpret_cast<const char *>(in), size).toHex();
-}
-
-
 xmrig::String xmrig::Buffer::toHex() const
 {
     if (m_size == 0) {
diff --git a/src/base/tools/Buffer.h b/src/base/tools/Buffer.h
index 28f92b9e..363e4697 100644
--- a/src/base/tools/Buffer.h
+++ b/src/base/tools/Buffer.h
@@ -35,14 +35,15 @@ namespace xmrig {
 class Buffer
 {
 public:
-    Buffer();
-    Buffer(Buffer &&other);
+    Buffer() = default;
+    Buffer(Buffer &&other) noexcept;
     Buffer(const Buffer &other);
     Buffer(const char *data, size_t size);
     Buffer(size_t size);
     ~Buffer();
 
 
+    inline bool isEmpty() const                     { return size() == 0; }
     inline bool isEqual(const Buffer &other) const  { return m_size == other.m_size && (m_size == 0 || memcmp(m_data, other.m_data, m_size) == 0); }
     inline char *data()                             { return m_data; }
     inline const char *data() const                 { return m_data; }
@@ -55,7 +56,7 @@ public:
 
     inline bool operator!=(const Buffer &other) const   { return !isEqual(other); }
     inline bool operator==(const Buffer &other) const   { return isEqual(other); }
-    inline Buffer &operator=(Buffer &&other)            { move(std::move(other)); return *this; }
+    inline Buffer &operator=(Buffer &&other) noexcept   { move(std::move(other)); return *this; }
     inline Buffer &operator=(const Buffer &other)       { from(other); return *this; }
 
 
@@ -67,12 +68,13 @@ public:
     inline static bool fromHex(const char *in, size_t size, uint8_t *out) { return fromHex(reinterpret_cast<const uint8_t *>(in), size, out); }
     inline static Buffer fromHex(const char *data)                        { return fromHex(data, strlen(data)); }
     inline static Buffer fromHex(const String &str)                       { return fromHex(str.data(), str.size()); }
+    inline static String toHex(const char *in, size_t size)               { return Buffer(in, size).toHex(); }
+    inline static String toHex(const uint8_t *in, size_t size)            { return Buffer(reinterpret_cast<const char *>(in), size).toHex(); }
     inline static void toHex(const char *in, size_t size, char *out)      { return toHex(reinterpret_cast<const uint8_t *>(in), size, reinterpret_cast<uint8_t *>(out)); }
     inline static void toHex(const uint8_t *in, size_t size, char *out)   { return toHex(in, size, reinterpret_cast<uint8_t *>(out)); }
 
     static bool fromHex(const uint8_t *in, size_t size, uint8_t *out);
     static Buffer fromHex(const char *data, size_t size);
-    static String toHex(const uint8_t *in, size_t size);
     static void toHex(const uint8_t *in, size_t size, uint8_t *out);
     String toHex() const;
 
@@ -80,8 +82,8 @@ private:
     void copy(const char *data, size_t size);
     void move(Buffer &&other);
 
-    char *m_data;
-    size_t m_size;
+    char *m_data    = nullptr;
+    size_t m_size   = 0;
 };
 
 
diff --git a/src/crypto/rx/Rx.cpp b/src/crypto/rx/Rx.cpp
index 81c39f67..76ab07c4 100644
--- a/src/crypto/rx/Rx.cpp
+++ b/src/crypto/rx/Rx.cpp
@@ -131,7 +131,7 @@ public:
 
     inline bool isNUMA() const                  { return m_numa; }
     inline const Algorithm &algorithm() const   { return m_algorithm; }
-    inline const uint8_t *seed() const          { return m_seed; }
+    inline const Buffer &seed() const           { return m_seed; }
     inline size_t count() const                 { return isNUMA() ? datasets.size() : 1; }
     inline void asyncSend()                     { m_ready++; if (m_ready == count()) { uv_async_send(m_async); } }
 
@@ -221,14 +221,13 @@ public:
         m_numa      = numa && Cpu::info()->nodes() > 1;
         m_hugePages = hugePages;
         m_listener  = listener;
-
-        memcpy(m_seed, job.seedHash(), sizeof(m_seed));
+        m_seed      = job.seed();
     }
 
 
     inline bool isReady(const Job &job)
     {
-        return m_ready == count() && m_algorithm == job.algorithm() && memcmp(m_seed, job.seedHash(), sizeof(m_seed)) == 0;
+        return m_ready == count() && m_algorithm == job.algorithm() && m_seed == job.seed();
     }
 
 
@@ -245,9 +244,9 @@ private:
     Algorithm m_algorithm;
     bool m_hugePages        = true;
     bool m_numa             = true;
+    Buffer m_seed;
     IRxListener *m_listener = nullptr;
     size_t m_ready          = 0;
-    uint8_t m_seed[32]{ 0 };
     uv_async_t *m_async;
 };
 
@@ -269,7 +268,7 @@ bool xmrig::Rx::init(const Job &job, int initThreads, bool hugePages, bool numa,
 
     d_ptr->setState(job, hugePages, numa, listener);
     const uint32_t threads = initThreads < 1 ? static_cast<uint32_t>(Cpu::info()->threads()) : static_cast<uint32_t>(initThreads);
-    const String buf       = Buffer::toHex(job.seedHash(), 8);
+    const String buf       = Buffer::toHex(job.seed().data(), 8);
 
     LOG_INFO("%s" MAGENTA_BOLD("init dataset%s") " algo " WHITE_BOLD("%s (") CYAN_BOLD("%u") WHITE_BOLD(" threads)") BLACK_BOLD(" seed %s..."),
              tag,
diff --git a/src/crypto/rx/RxAlgo.cpp b/src/crypto/rx/RxAlgo.cpp
index b306a4bc..23cda1f7 100644
--- a/src/crypto/rx/RxAlgo.cpp
+++ b/src/crypto/rx/RxAlgo.cpp
@@ -55,6 +55,46 @@ uint32_t xmrig::RxAlgo::version(Algorithm::Id algorithm)
 }
 
 
+uint32_t xmrig::RxAlgo::programCount(Algorithm::Id algorithm)
+{
+    switch (algorithm) {
+    case Algorithm::RX_0:
+        return RandomX_MoneroConfig.ProgramCount;
+
+    case Algorithm::RX_WOW:
+        return RandomX_WowneroConfig.ProgramCount;
+
+    case Algorithm::RX_LOKI:
+        return RandomX_LokiConfig.ProgramCount;
+
+    default:
+        break;
+    }
+
+    return 0;
+}
+
+
+uint32_t xmrig::RxAlgo::programIterations(Algorithm::Id algorithm)
+{
+    switch (algorithm) {
+    case Algorithm::RX_0:
+        return RandomX_MoneroConfig.ProgramIterations;
+
+    case Algorithm::RX_WOW:
+        return RandomX_WowneroConfig.ProgramIterations;
+
+    case Algorithm::RX_LOKI:
+        return RandomX_LokiConfig.ProgramIterations;
+
+    default:
+        break;
+    }
+
+    return 0;
+}
+
+
 uint32_t xmrig::RxAlgo::programSize(Algorithm::Id algorithm)
 {
     switch (algorithm) {
diff --git a/src/crypto/rx/RxAlgo.h b/src/crypto/rx/RxAlgo.h
index 859e488d..b0b40ed7 100644
--- a/src/crypto/rx/RxAlgo.h
+++ b/src/crypto/rx/RxAlgo.h
@@ -43,6 +43,8 @@ class RxAlgo
 {
 public:
     static Algorithm::Id apply(Algorithm::Id algorithm);
+    static uint32_t programCount(Algorithm::Id algorithm);
+    static uint32_t programIterations(Algorithm::Id algorithm);
     static uint32_t programSize(Algorithm::Id algorithm);
     static uint32_t version(Algorithm::Id algorithm);
 };
diff --git a/src/crypto/rx/RxCache.cpp b/src/crypto/rx/RxCache.cpp
index 92c366a2..89f0cb58 100644
--- a/src/crypto/rx/RxCache.cpp
+++ b/src/crypto/rx/RxCache.cpp
@@ -34,8 +34,7 @@ static_assert(RANDOMX_FLAG_LARGE_PAGES == 1, "RANDOMX_FLAG_LARGE_PAGES flag mism
 
 
 
-xmrig::RxCache::RxCache(bool hugePages) :
-    m_seed()
+xmrig::RxCache::RxCache(bool hugePages)
 {
     if (hugePages) {
         m_flags = RANDOMX_FLAG_JIT | RANDOMX_FLAG_LARGE_PAGES;
@@ -62,14 +61,14 @@ xmrig::RxCache::~RxCache()
 }
 
 
-bool xmrig::RxCache::init(const uint8_t *seed)
+bool xmrig::RxCache::init(const Buffer &seed)
 {
     if (isReady(seed)) {
         return false;
     }
 
-    memcpy(m_seed, seed, sizeof(m_seed));
-    randomx_init_cache(m_cache, m_seed, sizeof(m_seed));
+    m_seed = seed;
+    randomx_init_cache(m_cache, m_seed.data(), sizeof(m_seed));
 
     m_initCount++;
 
@@ -77,7 +76,7 @@ bool xmrig::RxCache::init(const uint8_t *seed)
 }
 
 
-bool xmrig::RxCache::isReady(const uint8_t *seed) const
+bool xmrig::RxCache::isReady(const Buffer &seed) const
 {
-    return m_initCount && memcmp(m_seed, seed, sizeof(m_seed)) == 0;
+    return !m_seed.isEmpty() && m_seed == seed;
 }
diff --git a/src/crypto/rx/RxCache.h b/src/crypto/rx/RxCache.h
index 12d173a2..2e61dbe4 100644
--- a/src/crypto/rx/RxCache.h
+++ b/src/crypto/rx/RxCache.h
@@ -31,6 +31,8 @@
 #include <cstdint>
 
 
+#include "base/tools/Buffer.h"
+#include "base/tools/Object.h"
 #include "crypto/randomx/configuration.h"
 
 
@@ -44,26 +46,28 @@ namespace xmrig
 class RxCache
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(RxCache)
+
     RxCache(bool hugePages = true);
     ~RxCache();
 
     inline bool isHugePages() const         { return m_flags & 1; }
     inline bool isJIT() const               { return m_flags & 8; }
-    inline const uint8_t *seed() const      { return m_seed; }
+    inline const Buffer &seed() const       { return m_seed; }
     inline randomx_cache *get() const       { return m_cache; }
     inline uint64_t initCount() const       { return m_initCount; }
 
-    bool init(const uint8_t *seed);
+    bool init(const Buffer &seed);
 
     static inline constexpr size_t maxSize() { return RANDOMX_CACHE_MAX_SIZE; }
 
 private:
-    bool isReady(const uint8_t *seed) const;
+    bool isReady(const Buffer &seed) const;
 
+    Buffer m_seed;
     int m_flags            = 0;
     randomx_cache *m_cache = nullptr;
     uint64_t m_initCount   = 0;
-    uint8_t m_seed[32];
 };
 
 
diff --git a/src/crypto/rx/RxDataset.cpp b/src/crypto/rx/RxDataset.cpp
index 3373abee..9115e9f2 100644
--- a/src/crypto/rx/RxDataset.cpp
+++ b/src/crypto/rx/RxDataset.cpp
@@ -64,7 +64,7 @@ xmrig::RxDataset::~RxDataset()
 }
 
 
-bool xmrig::RxDataset::init(const uint8_t *seed, uint32_t numThreads)
+bool xmrig::RxDataset::init(const Buffer &seed, uint32_t numThreads)
 {
     cache()->init(seed);
 
@@ -112,3 +112,9 @@ std::pair<size_t, size_t> xmrig::RxDataset::hugePages() const
 
     return { count, total };
 }
+
+
+void *xmrig::RxDataset::raw() const
+{
+    return m_dataset ? randomx_get_dataset_memory(m_dataset) : nullptr;
+}
diff --git a/src/crypto/rx/RxDataset.h b/src/crypto/rx/RxDataset.h
index 491827a3..e1b359a4 100644
--- a/src/crypto/rx/RxDataset.h
+++ b/src/crypto/rx/RxDataset.h
@@ -40,6 +40,7 @@ namespace xmrig
 {
 
 
+class Buffer;
 class RxCache;
 
 
@@ -54,9 +55,11 @@ public:
     inline bool isHugePages() const     { return m_flags & 1; }
     inline randomx_dataset *get() const { return m_dataset; }
     inline RxCache *cache() const       { return m_cache; }
+    inline size_t size() const          { return maxSize(); }
 
-    bool init(const uint8_t *seed, uint32_t numThreads);
+    bool init(const Buffer &seed, uint32_t numThreads);
     std::pair<size_t, size_t> hugePages() const;
+    void *raw() const;
 
     static inline constexpr size_t maxSize() { return RANDOMX_DATASET_MAX_SIZE; }
 
diff --git a/src/net/JobResults.cpp b/src/net/JobResults.cpp
index 7f4a56cb..4a6f87e9 100644
--- a/src/net/JobResults.cpp
+++ b/src/net/JobResults.cpp
@@ -23,17 +23,21 @@
  */
 
 
-#include <assert.h>
-#include <list>
-#include <mutex>
-#include <uv.h>
-
+#include "net/JobResults.h"
 
 #include "base/io/log/Log.h"
 #include "base/tools/Handle.h"
+#include "base/tools/Object.h"
 #include "net/interfaces/IJobResultListener.h"
 #include "net/JobResult.h"
-#include "net/JobResults.h"
+
+
+#ifdef XMRIG_ALGO_RANDOMX
+#   include "crypto/randomx/randomx.h"
+#   include "crypto/rx/Rx.h"
+#   include "crypto/rx/RxVm.h"
+#endif
+
 
 #if defined(XMRIG_FEATURE_OPENCL) || defined(XMRIG_FEATURE_CUDA)
 #   include "base/tools/Baton.h"
@@ -44,6 +48,14 @@
 #endif
 
 
+#include <cassert>
+#include <list>
+#include <mutex>
+#include <uv.h>
+
+#include "base/tools/Chrono.h"
+
+
 namespace xmrig {
 
 
@@ -80,14 +92,45 @@ public:
 };
 
 
+static inline void checkHash(const JobBundle &bundle, std::vector<JobResult> &results, uint32_t nonce, uint8_t hash[32], uint32_t &errors)
+{
+    if (*reinterpret_cast<uint64_t*>(hash + 24) < bundle.job.target()) {
+        results.emplace_back(bundle.job, nonce, hash);
+    }
+    else {
+        LOG_ERR("COMPUTE ERROR"); // TODO Extend information.
+        errors++;
+    }
+}
+
+
 static void getResults(JobBundle &bundle, std::vector<JobResult> &results, uint32_t &errors, bool hwAES)
 {
-    const Algorithm &algorithm  = bundle.job.algorithm();
-    VirtualMemory *memory       = new VirtualMemory(algorithm.l3(), false);
-    uint8_t hash[32];
+    const auto &algorithm = bundle.job.algorithm();
+    auto memory           = new VirtualMemory(algorithm.l3(), false);
+    uint8_t hash[32]{ 0 };
 
     if (algorithm.family() == Algorithm::RANDOM_X) {
-        errors += bundle.nonces.size(); // TODO RANDOM_X
+#       ifdef XMRIG_ALGO_RANDOMX
+        RxDataset *dataset = Rx::dataset(bundle.job, 0);
+        if (dataset == nullptr) {
+            errors += bundle.nonces.size();
+
+            return;
+        }
+
+        auto vm = new RxVm(dataset, memory->scratchpad(), !hwAES);
+
+        for (uint32_t nonce : bundle.nonces) {
+            *bundle.job.nonce() = nonce;
+
+            randomx_calculate_hash(vm->get(), bundle.job.blob(), bundle.job.size(), hash);
+
+            checkHash(bundle, results, nonce, hash, errors);
+        }
+
+        delete vm;
+#       endif
     }
     else if (algorithm.family() == Algorithm::ARGON2) {
         errors += bundle.nonces.size(); // TODO ARGON2
@@ -101,13 +144,7 @@ static void getResults(JobBundle &bundle, std::vector<JobResult> &results, uint3
 
             CnHash::fn(algorithm, hwAES ? CnHash::AV_SINGLE : CnHash::AV_SINGLE_SOFT, Assembly::NONE)(bundle.job.blob(), bundle.job.size(), hash, ctx, bundle.job.height());
 
-            if (*reinterpret_cast<uint64_t*>(hash + 24) < bundle.job.target()) {
-                results.push_back(JobResult(bundle.job, nonce, hash));
-            }
-            else {
-                LOG_ERR("COMPUTE ERROR"); // TODO Extend information.
-                errors++;
-            }
+            checkHash(bundle, results, nonce, hash, errors);
         }
     }
 
@@ -119,6 +156,8 @@ static void getResults(JobBundle &bundle, std::vector<JobResult> &results, uint3
 class JobResultsPrivate
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(JobResultsPrivate)
+
     inline JobResultsPrivate(IJobResultListener *listener, bool hwAES) :
         m_hwAES(hwAES),
         m_listener(listener)
@@ -149,7 +188,7 @@ public:
     inline void submit(const Job &job, uint32_t *results, size_t count)
     {
         std::lock_guard<std::mutex> lock(m_mutex);
-        m_bundles.push_back(JobBundle(job, results, count));
+        m_bundles.emplace_back(job, results, count);
 
         uv_async_send(m_async);
     }
@@ -171,7 +210,7 @@ private:
         m_results.swap(results);
         m_mutex.unlock();
 
-        for (auto result : results) {
+        for (const auto &result : results) {
             m_listener->onJobResult(result);
         }
 
@@ -179,20 +218,20 @@ private:
             return;
         }
 
-        JobBaton *baton = new JobBaton(std::move(bundles), m_listener, m_hwAES);
+        auto baton = new JobBaton(std::move(bundles), m_listener, m_hwAES);
 
         uv_queue_work(uv_default_loop(), &baton->req,
             [](uv_work_t *req) {
-                JobBaton *baton = static_cast<JobBaton*>(req->data);
+                auto baton = static_cast<JobBaton*>(req->data);
 
                 for (JobBundle &bundle : baton->bundles) {
                     getResults(bundle, baton->results, baton->errors, baton->hwAES);
                 }
             },
             [](uv_work_t *req, int) {
-                JobBaton *baton = static_cast<JobBaton*>(req->data);
+                auto baton = static_cast<JobBaton*>(req->data);
 
-                for (auto result : baton->results) {
+                for (const auto &result : baton->results) {
                     baton->listener->onJobResult(result);
                 }
 
@@ -209,7 +248,7 @@ private:
         m_results.swap(results);
         m_mutex.unlock();
 
-        for (auto result : results) {
+        for (const auto &result : results) {
             m_listener->onJobResult(result);
         }
     }
diff --git a/src/net/JobResults.h b/src/net/JobResults.h
index fddfd697..576e7219 100644
--- a/src/net/JobResults.h
+++ b/src/net/JobResults.h
@@ -27,7 +27,7 @@
 #define XMRIG_JOBRESULTS_H
 
 
-#include <stdint.h>
+#include <cstdint>
 
 
 namespace xmrig {