Implemented remaining kernels.

src/backend/opencl/kernels/Cn2Kernel.cpp

@@ -0,0 +1,59 @@
+/* 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/Cn2Kernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+xmrig::Cn2Kernel::Cn2Kernel(cl_program program) : OclKernel(program, "cn2")
+{
+}
+
+
+bool 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);
+}
+
+
+// __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)
+{
+    if (!setArg(0, sizeof(cl_mem), &scratchpads) || !setArg(1, sizeof(cl_mem), &states) || !setArg(6, sizeof(uint32_t), &threads)) {
+        return false;
+    }
+
+    for (uint32_t i = 0; i < branches.size(); ++i) {
+        if (!setArg(i + 2, sizeof(cl_mem), &branches[i])) {
+            return false;
+        }
+    }
+
+    return true;
+}

src/backend/opencl/kernels/Cn2Kernel.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_CN2KERNEL_H
+#define XMRIG_CN2KERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+class Cn2Kernel : public OclKernel
+{
+public:
+    Cn2Kernel(cl_program program);
+    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);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_CN2KERNEL_H */

src/backend/opencl/kernels/CnBranchKernel.cpp

@@ -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/>.
+ */
+
+
+#include "backend/opencl/kernels/CnBranchKernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+namespace xmrig {
+
+
+static const char *names[4] = { "Blake", "Groestl", "JH", "Skein" };
+
+
+} // namespace xmrig
+
+
+xmrig::CnBranchKernel::CnBranchKernel(size_t index, cl_program program) : OclKernel(program, names[index])
+{
+}
+
+
+bool 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);
+}
+
+
+// __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)
+{
+    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);
+}

src/backend/opencl/kernels/CnBranchKernel.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_CNBRANCHKERNEL_H
+#define XMRIG_CNBRANCHKERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+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);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_CNBRANCHKERNEL_H */

src/backend/opencl/opencl.cmake

@@ -8,6 +8,8 @@ if (WITH_OPENCL)
         src/backend/opencl/interfaces/IOclRunner.h
         src/backend/opencl/kernels/Cn0Kernel.h
         src/backend/opencl/kernels/Cn1Kernel.h
+        src/backend/opencl/kernels/Cn2Kernel.h
+        src/backend/opencl/kernels/CnBranchKernel.h
         src/backend/opencl/OclBackend.h
         src/backend/opencl/OclCache.h
         src/backend/opencl/OclConfig.h
@@ -30,6 +32,8 @@ if (WITH_OPENCL)
         src/backend/opencl/cl/OclSource.cpp
         src/backend/opencl/kernels/Cn0Kernel.cpp
         src/backend/opencl/kernels/Cn1Kernel.cpp
+        src/backend/opencl/kernels/Cn2Kernel.cpp
+        src/backend/opencl/kernels/CnBranchKernel.cpp
         src/backend/opencl/OclBackend.cpp
         src/backend/opencl/OclCache.cpp
         src/backend/opencl/OclConfig.cpp

src/backend/opencl/runners/OclCnRunner.cpp

@@ -25,6 +25,8 @@
 
 #include "backend/opencl/kernels/Cn0Kernel.h"
 #include "backend/opencl/kernels/Cn1Kernel.h"
+#include "backend/opencl/kernels/Cn2Kernel.h"
+#include "backend/opencl/kernels/CnBranchKernel.h"
 #include "backend/opencl/OclLaunchData.h"
 #include "backend/opencl/runners/OclCnRunner.h"
 #include "backend/opencl/wrappers/OclLib.h"
@@ -89,11 +91,14 @@ xmrig::OclCnRunner::OclCnRunner(size_t index, const OclLaunchData &data) : OclBa
 xmrig::OclCnRunner::~OclCnRunner()
 {
     delete m_cn0;
+    delete m_cn1;
+    delete m_cn2;
 
     OclLib::releaseMemObject(m_scratchpads);
     OclLib::releaseMemObject(m_states);
 
     for (size_t i = 0; i < BRANCH_MAX; ++i) {
+        delete m_branchKernels[i];
         OclLib::releaseMemObject(m_branches[i]);
     }
 }
@@ -115,9 +120,6 @@ bool xmrig::OclCnRunner::run(uint32_t nonce, uint32_t *hashOutput)
 {
     static const cl_uint zero = 0;
 
-    cl_int ret;
-    size_t branchNonces[4] = { 0 };
-
     const size_t g_intensity = data().thread.intensity();
     const size_t w_size      = data().thread.worksize();
     const size_t g_thd       = ((g_intensity + w_size - 1u) / w_size) * w_size;
@@ -142,7 +144,24 @@ bool xmrig::OclCnRunner::run(uint32_t nonce, uint32_t *hashOutput)
         return false;
     }
 
-    OclLib::finish(m_queue);
+    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;
+    }
+
+    uint32_t &results = hashOutput[0xFF];
+    if (results > 0xFF) {
+        results = 0xFF;
+    }
 
     return true;
 }
@@ -152,7 +171,7 @@ bool xmrig::OclCnRunner::selfTest() const
 {
     return OclBaseRunner::selfTest() &&
            m_cn0->isValid() &&
-           m_cn1->isValid();
+           m_cn1->isValid() && m_cn2->isValid();
 }
 
 
@@ -177,6 +196,18 @@ bool xmrig::OclCnRunner::set(const Job &job, uint8_t *blob)
         return false;
     }
 
+    const uint32_t intensity = data().thread.intensity();
+
+    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;
+        }
+    }
+
     return true;
 }
 
@@ -191,4 +222,9 @@ void xmrig::OclCnRunner::build()
 
     m_cn0 = new Cn0Kernel(m_program);
     m_cn1 = new Cn1Kernel(m_program);
+    m_cn2 = new Cn2Kernel(m_program);
+
+    for (size_t i = 0; i < BRANCH_MAX; ++i) {
+        m_branchKernels[i] = new CnBranchKernel(i, m_program);
+    }
 }

src/backend/opencl/runners/OclCnRunner.h

@@ -34,6 +34,8 @@ namespace xmrig {
 
 class Cn0Kernel;
 class Cn1Kernel;
+class Cn2Kernel;
+class CnBranchKernel;
 
 
 class OclCnRunner : public OclBaseRunner
@@ -59,11 +61,14 @@ private:
     };
 
 
-    cl_mem m_branches[BRANCH_MAX]   = { nullptr, nullptr, nullptr, nullptr };
+    cl_mem m_scratchpads    = nullptr;
-    cl_mem m_scratchpads            = nullptr;
+    cl_mem m_states         = nullptr;
-    cl_mem m_states                 = nullptr;
+    Cn0Kernel *m_cn0        = nullptr;
-    Cn0Kernel *m_cn0                = nullptr;
+    Cn1Kernel *m_cn1        = nullptr;
-    Cn1Kernel *m_cn1                = nullptr;
+    Cn2Kernel *m_cn2        = nullptr;
+
+    std::vector<cl_mem> m_branches                = { nullptr, nullptr, nullptr, nullptr };
+    std::vector<CnBranchKernel *> m_branchKernels = { nullptr, nullptr, nullptr, nullptr };
 };