Added Cn0Kernel and OclKernel classes.

src/backend/opencl/OclWorker.cpp

@@ -91,6 +91,8 @@ bool xmrig::OclWorker::selfTest()
 
 void xmrig::OclWorker::start()
 {
+    cl_uint results[0x100];
+
     while (Nonce::sequence(Nonce::OPENCL) > 0) {
         if (Nonce::isPaused()) {
             do {
@@ -116,6 +118,7 @@ void xmrig::OclWorker::start()
                 break;
             }
 
+            m_runner->run(results);
             std::this_thread::sleep_for(std::chrono::milliseconds(2000)); // FIXME
 
             m_job.nextRound(kReserveCount);

src/backend/opencl/kernels/Cn0Kernel.cpp

@@ -0,0 +1,42 @@
+/* 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/Cn0Kernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+xmrig::Cn0Kernel::Cn0Kernel(cl_program program) : OclKernel(program, "cn0")
+{
+}
+
+
+// __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(3, sizeof(uint32_t), &threads);
+}

src/backend/opencl/kernels/Cn0Kernel.h

@@ -0,0 +1,49 @@
+/* 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_CN0KERNEL_H
+#define XMRIG_CN0KERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+typedef struct _cl_mem *cl_mem;
+
+
+namespace xmrig {
+
+
+class Cn0Kernel : public OclKernel
+{
+public:
+    Cn0Kernel(cl_program program);
+    bool setArgs(cl_mem input, cl_mem scratchpads, cl_mem states, uint32_t threads);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_CN0KERNEL_H */

src/backend/opencl/opencl.cmake

@@ -6,6 +6,7 @@ if (WITH_OPENCL)
     set(HEADERS_BACKEND_OPENCL
         src/backend/opencl/cl/OclSource.h
         src/backend/opencl/interfaces/IOclRunner.h
+        src/backend/opencl/kernels/Cn0Kernel.h
         src/backend/opencl/OclBackend.h
         src/backend/opencl/OclCache.h
         src/backend/opencl/OclConfig.h
@@ -18,6 +19,7 @@ if (WITH_OPENCL)
         src/backend/opencl/wrappers/OclContext.h
         src/backend/opencl/wrappers/OclDevice.h
         src/backend/opencl/wrappers/OclError.h
+        src/backend/opencl/wrappers/OclKernel.h
         src/backend/opencl/wrappers/OclLib.h
         src/backend/opencl/wrappers/OclPlatform.h
         src/backend/opencl/wrappers/OclVendor.h
@@ -25,6 +27,7 @@ if (WITH_OPENCL)
 
     set(SOURCES_BACKEND_OPENCL
         src/backend/opencl/cl/OclSource.cpp
+        src/backend/opencl/kernels/Cn0Kernel.cpp
         src/backend/opencl/OclBackend.cpp
         src/backend/opencl/OclCache.cpp
         src/backend/opencl/OclConfig.cpp
@@ -37,6 +40,7 @@ if (WITH_OPENCL)
         src/backend/opencl/wrappers/OclContext.cpp
         src/backend/opencl/wrappers/OclDevice.cpp
         src/backend/opencl/wrappers/OclError.cpp
+        src/backend/opencl/wrappers/OclKernel.cpp
         src/backend/opencl/wrappers/OclLib.cpp
         src/backend/opencl/wrappers/OclPlatform.cpp
        )

src/backend/opencl/runners/OclCnRunner.cpp

@@ -23,6 +23,7 @@
  */
 
 
+#include "backend/opencl/kernels/Cn0Kernel.h"
 #include "backend/opencl/OclLaunchData.h"
 #include "backend/opencl/runners/OclCnRunner.h"
 #include "backend/opencl/wrappers/OclLib.h"
@@ -80,6 +81,8 @@ xmrig::OclCnRunner::OclCnRunner(size_t index, const OclLaunchData &data) : OclBa
 
 xmrig::OclCnRunner::~OclCnRunner()
 {
+    delete m_cn0;
+
     OclLib::releaseMemObject(m_scratchpads);
     OclLib::releaseMemObject(m_states);
     OclLib::releaseMemObject(m_blake256);
@@ -101,6 +104,12 @@ bool xmrig::OclCnRunner::isReadyToBuild() const
 }
 
 
+bool xmrig::OclCnRunner::selfTest() const
+{
+    return OclBaseRunner::selfTest() && m_cn0->isValid();
+}
+
+
 bool xmrig::OclCnRunner::set(const Job &job, uint8_t *blob)
 {
     if (job.size() > (Job::kMaxBlobSize - 4)) {
@@ -114,5 +123,22 @@ bool xmrig::OclCnRunner::set(const Job &job, uint8_t *blob)
         return false;
     }
 
+    if (!m_cn0->setArgs(m_input, m_scratchpads, m_states, data().thread.intensity())) {
+        return false;
+    }
+
+    LOG_WARN(GREEN_S "OK");
     return false;
 }
+
+
+void xmrig::OclCnRunner::build()
+{
+    OclBaseRunner::build();
+
+    if (!m_program) {
+        return;
+    }
+
+    m_cn0 = new Cn0Kernel(m_program);
+}

src/backend/opencl/runners/OclCnRunner.h

@@ -32,6 +32,9 @@
 namespace xmrig {
 
 
+class Cn0Kernel;
+
+
 class OclCnRunner : public OclBaseRunner
 {
 public:
@@ -40,7 +43,9 @@ public:
 
 protected:
     bool isReadyToBuild() const override;
+    bool selfTest() const override;
     bool set(const Job &job, uint8_t *blob) override;
+    void build() override;
 
 private:
     cl_mem m_blake256       = nullptr;
@@ -49,6 +54,7 @@ private:
     cl_mem m_scratchpads    = nullptr;
     cl_mem m_skein512       = nullptr;
     cl_mem m_states         = nullptr;
+    Cn0Kernel *m_cn0        = nullptr;
 };
 
 

src/backend/opencl/wrappers/OclKernel.cpp

@@ -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/>.
+ */
+
+
+#include "backend/opencl/wrappers/OclError.h"
+#include "backend/opencl/wrappers/OclKernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+#include "base/io/log/Log.h"
+
+
+xmrig::OclKernel::OclKernel(cl_program program, const char *name) :
+    m_name(name)
+{
+    cl_int ret = 0;
+    m_kernel = OclLib::createKernel(program, name, &ret);
+}
+
+
+xmrig::OclKernel::~OclKernel()
+{
+    OclLib::releaseKernel(m_kernel);
+}
+
+
+bool 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")
+                RED(" argument ") RED_BOLD("%u") RED(" size ") RED_BOLD("%zu"),
+                OclError::toString(ret), name().data(), index, size);
+    }
+
+    return ret == CL_SUCCESS;
+}

src/backend/opencl/wrappers/OclKernel.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_OCLKERNEL_H
+#define XMRIG_OCLKERNEL_H
+
+
+#include "base/tools/String.h"
+
+
+typedef struct _cl_kernel *cl_kernel;
+typedef struct _cl_program *cl_program;
+
+
+namespace xmrig {
+
+
+class OclKernel
+{
+public:
+    OclKernel(cl_program program, const char *name);
+    virtual ~OclKernel();
+
+    inline bool isValid() const         { return m_kernel != nullptr; }
+    inline cl_kernel kernel() const     { return m_kernel; }
+    inline const String &name() const   { return m_name; }
+
+    bool setArg(uint32_t index, size_t size, const void *value);
+
+private:
+    cl_kernel m_kernel = nullptr;
+    const String m_name;
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_OCLKERNEL_H */

src/backend/opencl/wrappers/OclLib.cpp

@@ -469,7 +469,10 @@ cl_kernel xmrig::OclLib::createKernel(cl_program program, const char *kernel_nam
 
     auto result = pCreateKernel(program, kernel_name, errcode_ret);
     if (*errcode_ret != CL_SUCCESS) {
-        LOG_ERR("Error %s when calling clCreateKernel for kernel %s.", OclError::toString(*errcode_ret), kernel_name);
+        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);
+
+        return nullptr;
     }
 
     return result;