Implemented verification on CPU.

src/backend/cpu/CpuWorker.cpp

@@ -208,7 +208,7 @@ void xmrig::CpuWorker<N>::start()
 
             for (size_t i = 0; i < N; ++i) {
                 if (*reinterpret_cast<uint64_t*>(m_hash + (i * 32) + 24) < job.target()) {
-                    JobResults::submit(JobResult(job, *m_job.nonce(i), m_hash + (i * 32)));
+                    JobResults::submit(job, *m_job.nonce(i), m_hash + (i * 32));
                 }
             }
 

src/backend/opencl/OclLaunchData.cpp

@@ -41,6 +41,6 @@ xmrig::OclLaunchData::OclLaunchData(const Miner *miner, const Algorithm &algorit
 
 bool xmrig::OclLaunchData::isEqual(const OclLaunchData &other) const
 {
-    return (other.algorithm.l3() == algorithm.l3() &&
-            other.thread         == thread);
+    return (other.algorithm == algorithm &&
+            other.thread    == thread);
 }

src/backend/opencl/OclWorker.cpp

@@ -68,7 +68,7 @@ xmrig::OclWorker::OclWorker(size_t id, const OclLaunchData &data) :
 
     case Algorithm::ARGON2:
 #       ifdef XMRIG_ALGO_ARGON2
-        m_runner = nullptr; // TODO
+        m_runner = nullptr; // TODO OclArgon2Runner
 #       endif
         break;
 
@@ -120,6 +120,10 @@ void xmrig::OclWorker::start()
                 return;
             }
 
+            if (results[0xFF] > 0) {
+                JobResults::submit(m_job.currentJob(), results, results[0xFF]);
+            }
+
             m_job.nextRound(roundSize(m_intensity), m_intensity);
             m_count += m_intensity;
 

src/backend/opencl/cl/cn/cryptonight.cl

@@ -74,7 +74,7 @@ inline ulong getIdx()
 
 
 __attribute__((reqd_work_group_size(8, 8, 1)))
-__kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states)
+__kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states, uint Threads)
 {
     uint ExpandedKey1[40];
     __local uint AES0[256], AES1[256], AES2[256], AES3[256];
@@ -476,7 +476,7 @@ __kernel void cn1_v2(__global uint4 *Scratchpad, __global ulong *states, uint va
 
 
 __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
-__kernel void cn1(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states)
+__kernel void cn1(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states, uint Threads)
 {
     ulong a[2], b[2];
     __local uint AES0[256], AES1[256];
@@ -548,11 +548,11 @@ __kernel void cn1(__global ulong *input, __global uint4 *Scratchpad, __global ul
                 long q = fast_div_heavy(n.s0, as_int4(n).s2 | 0x5);
                 *((__global long*)(Scratchpad + (IDX((idx0 & MASK) >> 4)))) = n.s0 ^ q;
 
-                if (variant == VARIANT_XHV) {
-                    idx0 = (~as_int4(n).s2) ^ q;
-                } else {
-                    idx0 = as_int4(n).s2 ^ q;
-                }
+#               if (ALGO == ALGO_CN_HEAVY_XHV) {
+                idx0 = (~as_int4(n).s2) ^ q;
+#               else
+                idx0 = as_int4(n).s2 ^ q;
+#               endif
             }
 #           endif
         }

src/backend/opencl/kernels/Cn0Kernel.cpp

@@ -42,10 +42,11 @@ bool xmrig::Cn0Kernel::enqueue(cl_command_queue queue, uint32_t nonce, size_t th
 }
 
 
-// __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states)
-bool xmrig::Cn0Kernel::setArgs(cl_mem input, cl_mem scratchpads, cl_mem states)
+// __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states, uint Threads)
+bool xmrig::Cn0Kernel::setArgs(cl_mem input, cl_mem scratchpads, cl_mem states, uint32_t threads)
 {
     return setArg(0, sizeof(cl_mem), &input) &&
            setArg(1, sizeof(cl_mem), &scratchpads) &&
-           setArg(2, sizeof(cl_mem), &states);
+           setArg(2, sizeof(cl_mem), &states) &&
+           setArg(3, sizeof(uint32_t), &threads);
 }

src/backend/opencl/kernels/Cn0Kernel.h

@@ -37,7 +37,7 @@ class Cn0Kernel : public OclKernel
 public:
     Cn0Kernel(cl_program program);
     bool enqueue(cl_command_queue queue, uint32_t nonce, size_t threads);
-    bool setArgs(cl_mem input, cl_mem scratchpads, cl_mem states);
+    bool setArgs(cl_mem input, cl_mem scratchpads, cl_mem states, uint32_t threads);
 };
 
 

src/backend/opencl/kernels/Cn1Kernel.cpp

@@ -42,10 +42,11 @@ bool xmrig::Cn1Kernel::enqueue(cl_command_queue queue, uint32_t nonce, size_t th
 }
 
 
-// __kernel void cn1(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states)
-bool xmrig::Cn1Kernel::setArgs(cl_mem input, cl_mem scratchpads, cl_mem states)
+// __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)
 {
     return setArg(0, sizeof(cl_mem), &input) &&
            setArg(1, sizeof(cl_mem), &scratchpads) &&
-           setArg(2, sizeof(cl_mem), &states);
+           setArg(2, sizeof(cl_mem), &states) &&
+           setArg(3, sizeof(uint32_t), &threads);
 }

src/backend/opencl/kernels/Cn1Kernel.h

@@ -37,7 +37,7 @@ class Cn1Kernel : public OclKernel
 public:
     Cn1Kernel(cl_program program);
     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);
+    bool setArgs(cl_mem input, cl_mem scratchpads, cl_mem states, uint32_t threads);
 };
 
 

src/backend/opencl/runners/OclCnRunner.cpp

@@ -188,16 +188,16 @@ bool xmrig::OclCnRunner::set(const Job &job, uint8_t *blob)
         return false;
     }
 
-    if (!m_cn0->setArgs(m_input, m_scratchpads, m_states)) {
-        return false;
-    }
-
-    if (!m_cn1->setArgs(m_input, m_scratchpads, m_states)) {
-        return false;
-    }
-
     const uint32_t intensity = data().thread.intensity();
 
+    if (!m_cn0->setArgs(m_input, m_scratchpads, m_states, intensity)) {
+        return false;
+    }
+
+    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;
     }

src/backend/opencl/runners/OclRxRunner.cpp

@@ -38,7 +38,7 @@ bool xmrig::OclRxRunner::run(uint32_t nonce, uint32_t *hashOutput)
 
 bool xmrig::OclRxRunner::selfTest() const
 {
-    return false; // TODO
+    return false; // TODO OclRxRunner
 }
 
 

src/net/JobResult.h

@@ -56,6 +56,7 @@ public:
 
     inline const uint8_t *result() const    { return m_result; }
     inline uint64_t actualDiff() const      { return Job::toDiff(reinterpret_cast<const uint64_t*>(m_result)[3]); }
+    inline uint8_t *result()                { return m_result; }
 
     const Algorithm algorithm;
     const String clientId;

src/net/JobResults.cpp

@@ -29,73 +29,205 @@
 #include <uv.h>
 
 
+#include "base/io/log/Log.h"
 #include "base/tools/Handle.h"
 #include "net/interfaces/IJobResultListener.h"
 #include "net/JobResult.h"
 #include "net/JobResults.h"
 
+#if defined(XMRIG_FEATURE_OPENCL) || defined(XMRIG_FEATURE_CUDA)
+#   include "base/tools/Baton.h"
+#   include "crypto/cn/CnCtx.h"
+#   include "crypto/cn/CnHash.h"
+#   include "crypto/cn/CryptoNight.h"
+#   include "crypto/common/VirtualMemory.h"
+#endif
+
+#include "base/tools/Buffer.h"
+
 
 namespace xmrig {
 
 
+#if defined(XMRIG_FEATURE_OPENCL) || defined(XMRIG_FEATURE_CUDA)
+class JobBundle
+{
+public:
+    inline JobBundle(const Job &job, uint32_t *results, size_t count) :
+        job(job),
+        nonces(count)
+    {
+        memcpy(nonces.data(), results, sizeof(uint32_t) * count);
+    }
+
+    Job job;
+    std::vector<uint32_t> nonces;
+};
+
+
+class JobBaton : public Baton<uv_work_t>
+{
+public:
+    inline JobBaton(std::list<JobBundle> &&bundles, IJobResultListener *listener, bool hwAES) :
+        hwAES(hwAES),
+        listener(listener),
+        bundles(std::move(bundles))
+    {}
+
+    const bool hwAES;
+    IJobResultListener *listener;
+    std::list<JobBundle> bundles;
+    std::vector<JobResult> results;
+    uint32_t errors = 0;
+};
+
+
+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];
+
+    if (algorithm.family() == Algorithm::RANDOM_X) {
+        errors += bundle.nonces.size(); // TODO RANDOM_X
+    }
+    else if (algorithm.family() == Algorithm::ARGON2) {
+        errors += bundle.nonces.size(); // TODO ARGON2
+    }
+    else {
+        cryptonight_ctx *ctx[1];
+        CnCtx::create(ctx, memory->scratchpad(), memory->size(), 1);
+
+        for (uint32_t nonce : bundle.nonces) {
+            *bundle.job.nonce() = nonce;
+
+            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++;
+            }
+        }
+    }
+
+    delete memory;
+}
+#endif
+
+
 class JobResultsPrivate
 {
 public:
-    inline JobResultsPrivate(IJobResultListener *listener) :
-        listener(listener)
+    inline JobResultsPrivate(IJobResultListener *listener, bool hwAES) :
+        m_hwAES(hwAES),
+        m_listener(listener)
     {
-        async = new uv_async_t;
-        async->data = this;
+        m_async = new uv_async_t;
+        m_async->data = this;
 
-        uv_async_init(uv_default_loop(), async, JobResultsPrivate::onResult);
+        uv_async_init(uv_default_loop(), m_async, JobResultsPrivate::onResult);
     }
 
 
     inline ~JobResultsPrivate()
     {
-        Handle::close(async);
+        Handle::close(m_async);
     }
 
 
-    void submit(const JobResult &result)
+    inline void submit(const JobResult &result)
     {
-        mutex.lock();
-        queue.push_back(result);
-        mutex.unlock();
+        std::lock_guard<std::mutex> lock(m_mutex);
+        m_results.push_back(result);
 
-        uv_async_send(async);
+        uv_async_send(m_async);
     }
 
 
+#   if defined(XMRIG_FEATURE_OPENCL) || defined(XMRIG_FEATURE_CUDA)
+    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));
+
+        uv_async_send(m_async);
+    }
+#   endif
+
+
 private:
     static void onResult(uv_async_t *handle) { static_cast<JobResultsPrivate*>(handle->data)->submit(); }
 
 
+#   if defined(XMRIG_FEATURE_OPENCL) || defined(XMRIG_FEATURE_CUDA)
+    inline void submit()
+    {
+        std::list<JobBundle> bundles;
+        std::list<JobResult> results;
+
+        m_mutex.lock();
+        m_bundles.swap(bundles);
+        m_results.swap(results);
+        m_mutex.unlock();
+
+        for (auto result : results) {
+            m_listener->onJobResult(result);
+        }
+
+        if (bundles.empty()) {
+            return;
+        }
+
+        JobBaton *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);
+
+                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);
+
+                for (auto result : baton->results) {
+                    baton->listener->onJobResult(result);
+                }
+
+                delete baton;
+            }
+        );
+    }
+#   else
     inline void submit()
     {
         std::list<JobResult> results;
 
-        mutex.lock();
-
-        while (!queue.empty()) {
-            results.push_back(std::move(queue.front()));
-            queue.pop_front();
-        }
-
-        mutex.unlock();
+        m_mutex.lock();
+        m_results.swap(results);
+        m_mutex.unlock();
 
         for (auto result : results) {
-            listener->onJobResult(result);
+            m_listener->onJobResult(result);
         }
-
-        results.clear();
     }
+#   endif
 
 
-    IJobResultListener *listener;
-    std::list<JobResult> queue;
-    std::mutex mutex;
-    uv_async_t *async;
+private:
+    const bool m_hwAES;
+    IJobResultListener *m_listener;
+    std::list<JobResult> m_results;
+    std::mutex m_mutex;
+    uv_async_t *m_async;
+
+#   if defined(XMRIG_FEATURE_OPENCL) || defined(XMRIG_FEATURE_CUDA)
+    std::list<JobBundle> m_bundles;
+#   endif
 };
 
 
@@ -106,11 +238,11 @@ static JobResultsPrivate *handler = nullptr;
 
 
 
-void xmrig::JobResults::setListener(IJobResultListener *listener)
+void xmrig::JobResults::setListener(IJobResultListener *listener, bool hwAES)
 {
     assert(handler == nullptr);
 
-    handler = new JobResultsPrivate(listener);
+    handler = new JobResultsPrivate(listener, hwAES);
 }
 
 
@@ -124,6 +256,12 @@ void xmrig::JobResults::stop()
 }
 
 
+void xmrig::JobResults::submit(const Job &job, uint32_t nonce, const uint8_t *result)
+{
+    submit(JobResult(job, nonce, result));
+}
+
+
 void xmrig::JobResults::submit(const JobResult &result)
 {
     assert(handler != nullptr);
@@ -132,3 +270,13 @@ void xmrig::JobResults::submit(const JobResult &result)
         handler->submit(result);
     }
 }
+
+
+#if defined(XMRIG_FEATURE_OPENCL) || defined(XMRIG_FEATURE_CUDA)
+void xmrig::JobResults::submit(const Job &job, uint32_t *results, size_t count)
+{
+    if (handler) {
+        handler->submit(job, results, count);
+    }
+}
+#endif

src/net/JobResults.h

@@ -27,19 +27,28 @@
 #define XMRIG_JOBRESULTS_H
 
 
+#include <stdint.h>
+
+
 namespace xmrig {
 
 
 class IJobResultListener;
+class Job;
 class JobResult;
 
 
 class JobResults
 {
 public:
-    static void setListener(IJobResultListener *listener);
+    static void setListener(IJobResultListener *listener, bool hwAES);
     static void stop();
+    static void submit(const Job &job, uint32_t nonce, const uint8_t *result);
     static void submit(const JobResult &result);
+
+#   if defined(XMRIG_FEATURE_OPENCL) || defined(XMRIG_FEATURE_CUDA)
+    static void submit(const Job &job, uint32_t *results, size_t count);
+#   endif
 };
 
 

src/net/Network.cpp

@@ -60,7 +60,7 @@ xmrig::Network::Network(Controller *controller) :
     m_donate(nullptr),
     m_timer(nullptr)
 {
-    JobResults::setListener(this);
+    JobResults::setListener(this, controller->config()->cpu().isHwAES());
     controller->addListener(this);
 
 #   ifdef XMRIG_FEATURE_API