Algo balancing

CMakeLists.txt

@@ -63,6 +63,7 @@ set(HEADERS
     src/net/strategies/DonateStrategy.h
     src/Summary.h
     src/version.h
+    src/workers/Benchmark.h
     src/workers/Handle.h
     src/workers/Hashrate.h
     src/workers/OclThread.h
@@ -120,6 +121,7 @@ set(SOURCES
     src/net/Network.cpp
     src/net/strategies/DonateStrategy.cpp
     src/Summary.cpp
+    src/workers/Benchmark.cpp
     src/workers/Handle.cpp
     src/workers/Hashrate.cpp
     src/workers/OclThread.cpp

build.bat

@@ -0,0 +1,18 @@
+@echo off
+call "C:\Program Files (x86)\Microsoft Visual Studio\2017\Community\VC\Auxiliary\Build\vcvars64.bat"
+rmdir /S /Q build
+del %~dp0\xmrig-amd-%1-win64.zip
+mkdir build &&^
+cd build &&^
+git clone https://github.com/MoneroOcean/xmrig-amd.git &&^
+git clone https://github.com/xmrig/xmrig-deps.git &&^
+mkdir xmrig-amd\build &&^
+cd xmrig-amd\build &&^
+git checkout %1 &&^
+cmake .. -G "Visual Studio 15 2017 Win64" -DXMRIG_DEPS=%~dp0\build\xmrig-deps\msvc2017\x64 &&^
+msbuild /p:Configuration=Release xmrig-amd.sln &&^
+cd Release &&^
+copy ..\..\src\config.json . &&^
+7za a -tzip -mx %~dp0\xmrig-amd-%1-win64.zip xmrig-amd.exe config.json &&^
+cd %~dp0 &&^
+rmdir /S /Q build

src/App.cpp

@@ -6,6 +6,7 @@
  * Copyright 2016      Jay D Dee   <[email protected]>
  * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
  * Copyright 2016-2018 XMRig       <https://github.com/xmrig>, <[email protected]>
+ * Copyright 2018 MoneroOcean      <https://github.com/MoneroOcean>, <[email protected]>
  *
  *   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
@@ -39,6 +40,7 @@
 #include "Summary.h"
 #include "version.h"
 #include "workers/Workers.h"
+#include "workers/Benchmark.h"
 
 
 #ifndef XMRIG_NO_HTTPD
@@ -83,6 +85,8 @@ App::~App()
 #   endif
 }
 
+// this should be global since we register onJobResult using this object method
+static Benchmark benchmark;
 
 int App::exec()
 {
@@ -128,7 +132,26 @@ int App::exec()
         return 1;
     }
 
-    m_controller->network()->connect();
+    // save config here to have option to store automatically generated "threads"
+    if (m_controller->config()->isShouldSave()) m_controller->config()->save();
+
+    // run benchmark before pool mining or not?
+    if (m_controller->config()->get_algo_perf(xmrig::PA_CN) == 0.0f || m_controller->config()->isCalibrateAlgo()) {
+        benchmark.set_controller(m_controller); // we need controller there to access config and network objects
+        benchmark.set_original_algorithm(m_controller->config()->algorithm());
+        Workers::setListener(&benchmark); // register benchmark as job reault listener to compute hashrates there
+        // write text before first benchmark round
+        Log::i()->text(m_controller->config()->isColors()
+            ? GREEN_BOLD(" >>>>> ") WHITE_BOLD("STARTING ALGO PERFORMANCE CALIBRATION (with %i seconds round)")
+            : " >>>>> STARTING ALGO PERFORMANCE CALIBRATION (with %i seconds round)",
+            m_controller->config()->calibrateAlgoTime()
+        );
+        // start benchmarking from first PerfAlgo in the list
+        if (m_controller->config()->get_algo_perf(xmrig::PA_CN) == 0.0f) benchmark.should_save_config();
+        benchmark.start_perf_bench(xmrig::PerfAlgo::PA_CN);
+    } else {
+        m_controller->network()->connect();
+    }
 
     const int r = uv_run(uv_default_loop(), UV_RUN_DEFAULT);
     uv_loop_close(uv_default_loop());

src/amd/GpuContext.h

@@ -6,6 +6,7 @@
  * Copyright 2016      Jay D Dee   <[email protected]>
  * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
  * Copyright 2016-2018 XMRig       <https://github.com/xmrig>, <[email protected]>
+ * Copyright 2018 MoneroOcean      <https://github.com/MoneroOcean>, <[email protected]>
  *
  *   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
@@ -25,7 +26,7 @@
 #define __GPUCONTEXT_H__
 
 
-#include "3rdparty/CL/cl.h"
+#include "amd/OclLib.h"
 
 
 #include <stdint.h>
@@ -73,6 +74,18 @@ struct GpuContext
         Nonce(0)
     {}
 
+    void release() { // stops all opencl kernels and releases all opencl resources
+        if (CommandQueues) {
+            OclLib::finish(CommandQueues);
+            OclLib::releaseCommandQueue(CommandQueues);
+        }
+        if (InputBuffer) OclLib::releaseMemObject(InputBuffer);
+        if (OutputBuffer) OclLib::releaseMemObject(OutputBuffer);
+        for (int i = 0; i < 6; ++ i)  if (ExtraBuffers[i]) OclLib::releaseMemObject(ExtraBuffers[i]);
+        for (int i = 0; i < 11; ++ i) if (Kernels[i]) OclLib::releaseKernel(Kernels[i]);
+        if (Program) OclLib::releaseProgram(Program);
+    }
+
     /*Input vars*/
     size_t deviceIdx;
     size_t rawIntensity;

src/amd/OclCLI.cpp

@@ -6,6 +6,7 @@
  * Copyright 2016      Jay D Dee   <[email protected]>
  * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
  * Copyright 2016-2018 XMRig       <https://github.com/xmrig>, <[email protected]>
+ * Copyright 2018 MoneroOcean      <https://github.com/MoneroOcean>, <[email protected]>
  *
  *   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
@@ -56,7 +57,7 @@ bool OclCLI::setup(std::vector<xmrig::IThread *> &threads)
 }
 
 
-void OclCLI::autoConf(std::vector<xmrig::IThread *> &threads, int *platformIndex, xmrig::Config *config)
+void OclCLI::autoConf(std::vector<xmrig::IThread *> &threads, int *platformIndex, const xmrig::Algorithm& algorithm, xmrig::Config *config)
 {
     *platformIndex = getAMDPlatformIdx(config);
     if (*platformIndex == -1) {
@@ -71,7 +72,7 @@ void OclCLI::autoConf(std::vector<xmrig::IThread *> &threads, int *platformIndex
     }
 
     constexpr size_t byteToMiB = 1024u * 1024u;
-    const size_t hashMemSize   = xmrig::cn_select_memory(config->algorithm().algo());
+    const size_t hashMemSize   = xmrig::cn_select_memory(algorithm.algo());
 
     for (GpuContext &ctx : devices) {
         // Vega APU slow and can cause BSOD, skip from autoconfig.
@@ -96,7 +97,7 @@ void OclCLI::autoConf(std::vector<xmrig::IThread *> &threads, int *platformIndex
             maxThreads = 2024u;
         }
 
-        if (config->algorithm().algo() == xmrig::CRYPTONIGHT_LITE) {
+        if (algorithm.algo() == xmrig::CRYPTONIGHT_LITE) {
             maxThreads *= 2u;
         }
 

src/amd/OclCLI.h

@@ -6,6 +6,7 @@
  * Copyright 2016      Jay D Dee   <[email protected]>
  * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
  * Copyright 2016-2018 XMRig       <https://github.com/xmrig>, <[email protected]>
+ * Copyright 2018 MoneroOcean      <https://github.com/MoneroOcean>, <[email protected]>
  *
  *   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
@@ -29,6 +30,7 @@
 
 
 #include "common/xmrig.h"
+#include "common/crypto/Algorithm.h" // need it for new xmrig::Algorithm autoConf param
 
 
 class OclThread;
@@ -46,7 +48,8 @@ class OclCLI
     OclCLI();
 
     bool setup(std::vector<xmrig::IThread *> &threads);
-    void autoConf(std::vector<xmrig::IThread *> &threads, int *platformIndex, xmrig::Config *config);
+    // autoConf now takes Algorithm parameter as input
+    void autoConf(std::vector<xmrig::IThread *> &threads, int *platformIndex, const xmrig::Algorithm&, xmrig::Config *config);
     void parseLaunch(const char *arg);
 
     inline void parseAffinity(const char *arg) { parse(m_affinity, arg); }

src/amd/OclGPU.cpp

@@ -7,6 +7,7 @@
  * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
  * Copyright 2018      Lee Clagett <https://github.com/vtnerd>
  * Copyright 2016-2018 XMRig       <https://github.com/xmrig>, <[email protected]>
+ * Copyright 2018 MoneroOcean      <https://github.com/MoneroOcean>, <[email protected]>
  *
  *   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
@@ -402,15 +403,18 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, xmrig::Config *config)
 #   ifdef __GNUC__
     cl_device_id TempDeviceList[num_gpus];
 #   else
-    cl_device_id* TempDeviceList = (cl_device_id*)_alloca(entries * sizeof(cl_device_id));
+    cl_device_id* TempDeviceList = (cl_device_id*)_alloca(num_gpus * sizeof(cl_device_id));
 #   endif
 
     for (size_t i = 0; i < num_gpus; ++i) {
         ctx[i].DeviceID = DeviceIDList[ctx[i].deviceIdx];
         TempDeviceList[i] = DeviceIDList[ctx[i].deviceIdx];
     }
 
-    cl_context opencl_ctx = OclLib::createContext(nullptr, num_gpus, TempDeviceList, nullptr, nullptr, &ret);
+    // we store previous OpenCL context in static variable to be able to release it next time we do algo switch
+    static cl_context opencl_ctx = nullptr;
+    if (opencl_ctx) OclLib::releaseContext(opencl_ctx);
+    opencl_ctx = OclLib::createContext(nullptr, num_gpus, TempDeviceList, nullptr, nullptr, &ret);
     if(ret != CL_SUCCESS) {
         return OCL_ERR_API;
     }

src/amd/OclLib.cpp

@@ -6,6 +6,7 @@
  * Copyright 2016      Jay D Dee   <[email protected]>
  * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
  * Copyright 2016-2018 XMRig       <https://github.com/xmrig>, <[email protected]>
+ * Copyright 2018 MoneroOcean      <https://github.com/MoneroOcean>, <[email protected]>
  *
  *   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
@@ -53,6 +54,11 @@ static const char *kGetPlatformInfo                  = "clGetPlatformInfo";
 static const char *kGetProgramBuildInfo              = "clGetProgramBuildInfo";
 static const char *kGetProgramInfo                   = "clGetProgramInfo";
 static const char *kSetKernelArg                     = "clSetKernelArg";
+static const char *kReleaseContext                   = "clReleaseContext";
+static const char *kReleaseProgram                   = "clReleaseProgram";
+static const char *kReleaseCommandQueue              = "clReleaseCommandQueue";
+static const char *kReleaseMemObject                 = "clReleaseMemObject";
+static const char *kReleaseKernel                    = "clReleaseKernel";
 
 typedef cl_command_queue (CL_API_CALL *createCommandQueueWithProperties_t)(cl_context, cl_device_id, const cl_queue_properties *, cl_int *);
 typedef cl_command_queue (CL_API_CALL *createCommandQueue_t)(cl_context, cl_device_id, cl_command_queue_properties, cl_int *);
@@ -73,6 +79,11 @@ typedef cl_kernel (CL_API_CALL *createKernel_t)(cl_program, const char *, cl_int
 typedef cl_mem (CL_API_CALL *createBuffer_t)(cl_context, cl_mem_flags, size_t, void *, cl_int *);
 typedef cl_program (CL_API_CALL *createProgramWithBinary_t)(cl_context, cl_uint, const cl_device_id *, const size_t *, const unsigned char **, cl_int *, cl_int *);
 typedef cl_program (CL_API_CALL *createProgramWithSource_t)(cl_context, cl_uint, const char **, const size_t *, cl_int *);
+typedef cl_int (CL_API_CALL *releaseContext_t)(cl_context);
+typedef cl_int (CL_API_CALL *releaseProgram_t)(cl_program);
+typedef cl_int (CL_API_CALL *releaseCommandQueue_t)(cl_command_queue);
+typedef cl_int (CL_API_CALL *releaseMemObject_t)(cl_mem);
+typedef cl_int (CL_API_CALL *releaseKernel_t)(cl_kernel);
 
 static createCommandQueueWithProperties_t pCreateCommandQueueWithProperties = nullptr;
 static createCommandQueue_t pCreateCommandQueue                             = nullptr;
@@ -93,6 +104,11 @@ static createKernel_t pCreateKernel                                         = nu
 static createBuffer_t pCreateBuffer                                         = nullptr;
 static createProgramWithBinary_t pCreateProgramWithBinary                   = nullptr;
 static createProgramWithSource_t pCreateProgramWithSource                   = nullptr;
+static releaseContext_t pReleaseContext                                     = nullptr;
+static releaseProgram_t pReleaseProgram                                     = nullptr;
+static releaseCommandQueue_t pReleaseCommandQueue                           = nullptr;
+static releaseMemObject_t pReleaseMemObject                                 = nullptr;
+static releaseKernel_t pReleaseKernel                                       = nullptr;
 
 #define DLSYM(x) if (uv_dlsym(&oclLib, k##x, reinterpret_cast<void**>(&p##x)) == -1) { return false; }
 
@@ -128,6 +144,11 @@ bool OclLib::load()
     DLSYM(CreateBuffer);
     DLSYM(CreateProgramWithBinary);
     DLSYM(CreateProgramWithSource);
+    DLSYM(ReleaseContext);
+    DLSYM(ReleaseProgram);
+    DLSYM(ReleaseCommandQueue);
+    DLSYM(ReleaseMemObject);
+    DLSYM(ReleaseKernel);
 
     uv_dlsym(&oclLib, kCreateCommandQueueWithProperties, reinterpret_cast<void**>(&pCreateCommandQueueWithProperties));
 
@@ -335,3 +356,42 @@ cl_program OclLib::createProgramWithSource(cl_context context, cl_uint count, co
 
     return result;
 }
+
+cl_int OclLib::releaseContext(cl_context context)
+{
+    assert(pReleaseContext != nullptr);
+
+    return pReleaseContext(context);
+}
+
+
+cl_int OclLib::releaseProgram(cl_program program)
+{
+    assert(pReleaseProgram != nullptr);
+
+    return pReleaseProgram(program);
+}
+
+
+cl_int OclLib::releaseCommandQueue(cl_command_queue command_queue)
+{
+    assert(pReleaseCommandQueue != nullptr);
+
+    return pReleaseCommandQueue(command_queue);
+}
+
+
+cl_int OclLib::releaseMemObject(cl_mem memobj)
+{
+    assert(pReleaseMemObject != nullptr);
+
+    return pReleaseMemObject(memobj);
+}
+
+
+cl_int OclLib::releaseKernel(cl_kernel kernel)
+{
+    assert(pReleaseKernel != nullptr);
+
+    return pReleaseKernel(kernel);
+}

src/amd/OclLib.h

@@ -6,6 +6,7 @@
  * Copyright 2016      Jay D Dee   <[email protected]>
  * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
  * Copyright 2016-2018 XMRig       <https://github.com/xmrig>, <[email protected]>
+ * Copyright 2018 MoneroOcean      <https://github.com/MoneroOcean>, <[email protected]>
  *
  *   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
@@ -51,6 +52,12 @@ class OclLib
     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);
+    // we need to properly release OpenCL we created to be able to do algo switching
+    static cl_int releaseContext(cl_context context);
+    static cl_int releaseProgram(cl_program program);
+    static cl_int releaseCommandQueue(cl_command_queue command_queue);
+    static cl_int releaseMemObject(cl_mem memobj);
+    static cl_int releaseKernel(cl_kernel kernel);
 
 private:
     static bool load();

src/amd/opencl/cryptonight.cl

@@ -468,9 +468,11 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul
     mem_fence(CLK_LOCAL_MEM_FENCE);
 
 #   if (ALGO == CRYPTONIGHT_HEAVY)
-    {
-        __local uint4 xin[8][WORKSIZE];
+    __local uint4 xin[8][WORKSIZE];
+#   endif
 
+#   if (ALGO == CRYPTONIGHT_HEAVY)
+    {
         /* Also left over threads perform this loop.
          * The left over thread results will be ignored
          */