Merge pull request #27 from Spudz76/ocean-master-fixups

Cleanup patchset

cmake/randomx.cmake

@@ -40,22 +40,22 @@ if (WITH_RANDOMX)
         src/crypto/rx/RxDataset.cpp
         src/crypto/rx/RxQueue.cpp
         src/crypto/rx/RxVm.cpp
-        src/crypto/defyx/align.h
+        src/crypto/randomx/defyx/align.h
-        src/crypto/defyx/brg_endian.h
+        src/crypto/randomx/defyx/brg_endian.h
-        src/crypto/defyx/KangarooTwelve.c
+        src/crypto/randomx/defyx/KangarooTwelve.c
-        src/crypto/defyx/KangarooTwelve.h
+        src/crypto/randomx/defyx/KangarooTwelve.h
-        src/crypto/defyx/KeccakP-1600-reference.c
+        src/crypto/randomx/defyx/KeccakP-1600-reference.c
-        src/crypto/defyx/KeccakP-1600-SnP.h
+        src/crypto/randomx/defyx/KeccakP-1600-SnP.h
-        src/crypto/defyx/KeccakSponge-common.h
+        src/crypto/randomx/defyx/KeccakSponge-common.h
-        src/crypto/defyx/KeccakSponge.inc
+        src/crypto/randomx/defyx/KeccakSponge.inc
-        src/crypto/defyx/KeccakSpongeWidth1600.c
+        src/crypto/randomx/defyx/KeccakSpongeWidth1600.c
-        src/crypto/defyx/KeccakSpongeWidth1600.h
+        src/crypto/randomx/defyx/KeccakSpongeWidth1600.h
-        src/crypto/defyx/Phases.h
+        src/crypto/randomx/defyx/Phases.h
-        src/crypto/defyx/sha256.c
+        src/crypto/randomx/defyx/sha256.c
-        src/crypto/defyx/sha256.h
+        src/crypto/randomx/defyx/sha256.h
-        src/crypto/defyx/sysendian.h
+        src/crypto/randomx/defyx/sysendian.h
-        src/crypto/defyx/yescrypt.h
+        src/crypto/randomx/defyx/yescrypt.h
-        src/crypto/defyx/yescrypt-best.c
+        src/crypto/randomx/defyx/yescrypt-best.c
     )
 
     if (CMAKE_C_COMPILER_ID MATCHES MSVC)

scripts/generate_cl.js

@@ -43,15 +43,6 @@ function cn_r()
 }
 
 
-function cn_gpu()
-{
-    const cn_gpu = opencl_minify(addIncludes('cryptonight_gpu.cl', [ 'wolf-aes.cl', 'keccak.cl' ]));
-
-    // fs.writeFileSync('cryptonight_gpu_gen.cl', cn_gpu);
-    fs.writeFileSync('cryptonight_gpu_cl.h', text2h(cn_gpu, 'xmrig', 'cryptonight_gpu_cl'));
-}
-
-
 function rx()
 {
     let rx = addIncludes('randomx.cl', [
@@ -96,6 +87,15 @@ function kawpow()
 }
 
 
+function cn_gpu()
+{
+    const cn_gpu = opencl_minify(addIncludes('cryptonight_gpu.cl', [ 'wolf-aes.cl', 'keccak.cl' ]));
+
+    // fs.writeFileSync('cryptonight_gpu_gen.cl', cn_gpu);
+    fs.writeFileSync('cryptonight_gpu_cl.h', text2h(cn_gpu, 'xmrig', 'cryptonight_gpu_cl'));
+}
+
+
 process.chdir(path.resolve('src/backend/opencl/cl/cn'));
 
 cn();

src/backend/cpu/CpuWorker.cpp

@@ -143,13 +143,13 @@ bool xmrig::CpuWorker<N>::selfTest()
                         verify(Algorithm::CN_XAO,    test_output_xao)  &&
                         verify(Algorithm::CN_RTO,    test_output_rto)  &&
                         verify(Algorithm::CN_HALF,   test_output_half) &&
-#                       ifdef XMRIG_ALGO_CN_GPU
-                        verify(Algorithm::CN_GPU,    test_output_gpu)  &&
-#                       endif
                         verify2(Algorithm::CN_R,     test_output_r)    &&
                         verify(Algorithm::CN_RWZ,    test_output_rwz)  &&
                         verify(Algorithm::CN_ZLS,    test_output_zls)  &&
                         verify(Algorithm::CN_CCX,    test_output_ccx)  &&
+#                       ifdef XMRIG_ALGO_CN_GPU
+                        verify(Algorithm::CN_GPU,    test_output_gpu)  &&
+#                       endif
                         verify(Algorithm::CN_DOUBLE, test_output_double);
 
 #       ifdef XMRIG_ALGO_CN_GPU
@@ -254,22 +254,14 @@ void xmrig::CpuWorker<N>::start()
 
                 if (first) {
                     first = false;
-                    if (job.algorithm() == Algorithm::RX_DEFYX) {
+                    randomx_calculate_hash_first(m_vm, tempHash, m_job.blob(), job.size(), job.algorithm());
-                        defyx_calculate_hash_first(m_vm, tempHash, m_job.blob(), job.size());
-                    } else {
-                        randomx_calculate_hash_first(m_vm, tempHash, m_job.blob(), job.size());
-                    }
                 }
 
                 if (!nextRound(m_job)) {
                     break;
                 }
 
-                if (job.algorithm() == Algorithm::RX_DEFYX) {
+                randomx_calculate_hash_next(m_vm, tempHash, m_job.blob(), job.size(), m_hash, job.algorithm());
-                    defyx_calculate_hash_next(m_vm, tempHash, m_job.blob(), job.size(), m_hash);
-                } else {
-                    randomx_calculate_hash_next(m_vm, tempHash, m_job.blob(), job.size(), m_hash);
-                }
             }
             else
 #           endif

src/backend/cpu/platform/BasicCpuInfo.cpp

@@ -228,12 +228,6 @@ xmrig::CpuThreads xmrig::BasicCpuInfo::threads(const Algorithm &algorithm, uint3
         return 1;
     }
 
-#   ifdef XMRIG_ALGO_CN_GPU
-    if (algorithm == Algorithm::CN_GPU) {
-        return count;
-    }
-#   endif
-
 #   ifdef XMRIG_ALGO_CN_LITE
     if (algorithm.family() == Algorithm::CN_LITE) {
         return CpuThreads(count, 1);
@@ -278,6 +272,12 @@ xmrig::CpuThreads xmrig::BasicCpuInfo::threads(const Algorithm &algorithm, uint3
     }
 #   endif
 
+#   ifdef XMRIG_ALGO_CN_GPU
+    if (algorithm == Algorithm::CN_GPU) {
+        return count;
+    }
+#   endif
+
     return CpuThreads(std::max<size_t>(count / 2, 1), 1);
 }
 

src/backend/opencl/OclThread.h

@@ -51,20 +51,6 @@ public:
         setIntensity(intensity);
     }
 
-#   ifdef XMRIG_ALGO_CN_GPU
-    OclThread(uint32_t index, uint32_t intensity, uint32_t worksize, uint32_t threads, uint32_t unrollFactor) :
-        m_fields(0),
-        m_threads(threads, -1),
-        m_index(index),
-        m_memChunk(0),
-        m_stridedIndex(0),
-        m_unrollFactor(unrollFactor),
-        m_worksize(worksize)
-    {
-        setIntensity(intensity);
-    }
-#   endif
-
 #   ifdef XMRIG_ALGO_RANDOMX
     OclThread(uint32_t index, uint32_t intensity, uint32_t worksize, uint32_t threads, bool gcnAsm, bool datasetHost, uint32_t bfactor) :
         m_datasetHost(datasetHost),
@@ -109,6 +95,20 @@ public:
     }
 #   endif
 
+#   ifdef XMRIG_ALGO_CN_GPU
+    OclThread(uint32_t index, uint32_t intensity, uint32_t worksize, uint32_t threads, uint32_t unrollFactor) :
+        m_fields(0),
+        m_threads(threads, -1),
+        m_index(index),
+        m_memChunk(0),
+        m_stridedIndex(0),
+        m_unrollFactor(unrollFactor),
+        m_worksize(worksize)
+    {
+        setIntensity(intensity);
+    }
+#   endif
+
     OclThread(const rapidjson::Value &value);
 
     inline bool isAsm() const                               { return m_gcnAsm; }

src/backend/opencl/OclWorker.cpp

@@ -116,14 +116,16 @@ xmrig::OclWorker::OclWorker(size_t id, const OclLaunchData &data) :
 #       endif
         break;
 
-    case Algorithm::CN_GPU:
-#       ifdef XMRIG_ALGO_CN_GPU
-        m_runner = new OclRyoRunner(id, data);
-#       endif
-        break;
-
     default:
-        m_runner = new OclCnRunner(id, data);
+#       ifdef XMRIG_ALGO_CN_GPU
+        if (m_algorithm == Algorithm::CN_GPU) {
+            m_runner = new OclRyoRunner(id, data);
+        }
+        else
+#       endif
+        {
+            m_runner = new OclCnRunner(id, data);
+        }
         break;
     }
 

src/backend/opencl/cl/OclSource.cpp

@@ -28,10 +28,6 @@
 #include "base/crypto/Algorithm.h"
 
 
-#ifdef XMRIG_ALGO_CN_GPU
-#   include "backend/opencl/cl/cn/cryptonight_gpu_cl.h"
-#endif
-
 #ifdef XMRIG_ALGO_RANDOMX
 #   include "backend/opencl/cl/rx/randomx_cl.h"
 #endif
@@ -45,6 +41,10 @@
 #   include "backend/opencl/cl/kawpow/kawpow_dag_cl.h"
 #endif
 
+#ifdef XMRIG_ALGO_CN_GPU
+#   include "backend/opencl/cl/cn/cryptonight_gpu_cl.h"
+#endif
+
 
 const char *xmrig::OclSource::get(const Algorithm &algorithm)
 {

src/backend/opencl/cl/cn/algorithm.cl

@@ -9,26 +9,26 @@
 #define ALGO_CN_RWZ         8
 #define ALGO_CN_ZLS         9
 #define ALGO_CN_DOUBLE      10
-#define ALGO_CN_GPU         11
+#define ALGO_CN_LITE_0      11
-#define ALGO_CN_LITE_0      12
+#define ALGO_CN_LITE_1      12
-#define ALGO_CN_LITE_1      13
+#define ALGO_CN_HEAVY_0     13
-#define ALGO_CN_HEAVY_0     14
+#define ALGO_CN_HEAVY_TUBE  14
-#define ALGO_CN_HEAVY_TUBE  15
+#define ALGO_CN_HEAVY_XHV   15
-#define ALGO_CN_HEAVY_XHV   16
+#define ALGO_CN_PICO_0      16
-#define ALGO_CN_PICO_0      17
+#define ALGO_CN_PICO_TLO    17
-#define ALGO_CN_PICO_TLO    18
+#define ALGO_CN_CCX         18
-#define ALGO_CN_CCX         19
+#define ALGO_CN_GPU         19
 #define ALGO_RX_0           20
 #define ALGO_RX_WOW         21
 #define ALGO_RX_LOKI        22
 #define ALGO_RX_ARQ         23
 #define ALGO_RX_SFX         24
 #define ALGO_RX_KEVA        25
-#define ALGO_RX_DEFYX       26
+#define ALGO_AR2_CHUKWA     26
-#define ALGO_AR2_CHUKWA     27
+#define ALGO_AR2_WRKZ       27
-#define ALGO_AR2_WRKZ       28
+#define ALGO_ASTROBWT_DERO  28
-#define ALGO_ASTROBWT_DERO  29
+#define ALGO_KAWPOW_RVN     29
-#define ALGO_KAWPOW_RVN     30
+#define ALGO_RX_DEFYX       30
 
 #define FAMILY_UNKNOWN      0
 #define FAMILY_CN           1

src/backend/opencl/cl/cn/cryptonight_gpu.cl

@@ -87,7 +87,7 @@ inline global int4* scratchpad_ptr(uint idx, uint n, __global int *lpad) { retur
 
 inline float4 fma_break(float4 x)
 {
-    // Break the dependency chain by setitng the exp to ?????01
+    // Break the dependency chain by setting the exp to ?????01
     x = _mm_and_ps(x, 0xFEFFFFFF);
     return _mm_or_ps(x, 0x00800000);
 }

src/backend/opencl/cl/rx/randomx_cl.h

@@ -11,22 +11,22 @@ static const char randomx_cl[130649] = {
     0x5f,0x43,0x4e,0x5f,0x52,0x54,0x4f,0x20,0x37,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x43,0x4e,0x5f,0x52,0x57,0x5a,0x20,0x38,0x0a,
     0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x43,0x4e,0x5f,0x5a,0x4c,0x53,0x20,0x39,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,
     0x47,0x4f,0x5f,0x43,0x4e,0x5f,0x44,0x4f,0x55,0x42,0x4c,0x45,0x20,0x31,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x43,0x4e,0x5f,
-    0x47,0x50,0x55,0x20,0x31,0x31,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x43,0x4e,0x5f,0x4c,0x49,0x54,0x45,0x5f,0x30,0x20,0x31,0x32,
+    0x4c,0x49,0x54,0x45,0x5f,0x30,0x20,0x31,0x31,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x43,0x4e,0x5f,0x4c,0x49,0x54,0x45,0x5f,0x31,
-    0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x43,0x4e,0x5f,0x4c,0x49,0x54,0x45,0x5f,0x31,0x20,0x31,0x33,0x0a,0x23,0x64,0x65,0x66,0x69,
+    0x20,0x31,0x32,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x43,0x4e,0x5f,0x48,0x45,0x41,0x56,0x59,0x5f,0x30,0x20,0x31,0x33,0x0a,0x23,
-    0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x43,0x4e,0x5f,0x48,0x45,0x41,0x56,0x59,0x5f,0x30,0x20,0x31,0x34,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,
+    0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x43,0x4e,0x5f,0x48,0x45,0x41,0x56,0x59,0x5f,0x54,0x55,0x42,0x45,0x20,0x31,0x34,0x0a,0x23,0x64,0x65,
-    0x47,0x4f,0x5f,0x43,0x4e,0x5f,0x48,0x45,0x41,0x56,0x59,0x5f,0x54,0x55,0x42,0x45,0x20,0x31,0x35,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,
+    0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x43,0x4e,0x5f,0x48,0x45,0x41,0x56,0x59,0x5f,0x58,0x48,0x56,0x20,0x31,0x35,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,
-    0x5f,0x43,0x4e,0x5f,0x48,0x45,0x41,0x56,0x59,0x5f,0x58,0x48,0x56,0x20,0x31,0x36,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x43,0x4e,
+    0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x43,0x4e,0x5f,0x50,0x49,0x43,0x4f,0x5f,0x30,0x20,0x31,0x36,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,
-    0x5f,0x50,0x49,0x43,0x4f,0x5f,0x30,0x20,0x31,0x37,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x43,0x4e,0x5f,0x50,0x49,0x43,0x4f,0x5f,
+    0x5f,0x43,0x4e,0x5f,0x50,0x49,0x43,0x4f,0x5f,0x54,0x4c,0x4f,0x20,0x31,0x37,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x43,0x4e,0x5f,
-    0x54,0x4c,0x4f,0x20,0x31,0x38,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x43,0x4e,0x5f,0x43,0x43,0x58,0x20,0x31,0x39,0x0a,0x23,0x64,
+    0x43,0x43,0x58,0x20,0x31,0x38,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x43,0x4e,0x5f,0x47,0x50,0x55,0x20,0x31,0x39,0x0a,0x23,0x64,
     0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x30,0x20,0x32,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,
     0x52,0x58,0x5f,0x57,0x4f,0x57,0x20,0x32,0x31,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x4c,0x4f,0x4b,0x49,0x20,0x32,
     0x32,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x41,0x52,0x51,0x20,0x32,0x33,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,
     0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x53,0x46,0x58,0x20,0x32,0x34,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,
-    0x4b,0x45,0x56,0x41,0x20,0x32,0x35,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x44,0x45,0x46,0x59,0x58,0x20,0x32,0x36,
+    0x4b,0x45,0x56,0x41,0x20,0x32,0x35,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,0x41,0x20,
-    0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,0x41,0x20,0x32,0x37,0x0a,0x23,0x64,0x65,0x66,
+    0x32,0x36,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x57,0x52,0x4b,0x5a,0x20,0x32,0x37,0x0a,0x23,0x64,0x65,0x66,
-    0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x57,0x52,0x4b,0x5a,0x20,0x32,0x38,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,
+    0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x53,0x54,0x52,0x4f,0x42,0x57,0x54,0x5f,0x44,0x45,0x52,0x4f,0x20,0x32,0x38,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,
-    0x4f,0x5f,0x41,0x53,0x54,0x52,0x4f,0x42,0x57,0x54,0x5f,0x44,0x45,0x52,0x4f,0x20,0x32,0x39,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,
+    0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x4b,0x41,0x57,0x50,0x4f,0x57,0x5f,0x52,0x56,0x4e,0x20,0x32,0x39,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,
-    0x4b,0x41,0x57,0x50,0x4f,0x57,0x5f,0x52,0x56,0x4e,0x20,0x33,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x41,0x4d,0x49,0x4c,0x59,0x5f,0x55,0x4e,0x4b,
+    0x4f,0x5f,0x52,0x58,0x5f,0x44,0x45,0x46,0x59,0x58,0x20,0x33,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x41,0x4d,0x49,0x4c,0x59,0x5f,0x55,0x4e,0x4b,
     0x4e,0x4f,0x57,0x4e,0x20,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x41,0x4d,0x49,0x4c,0x59,0x5f,0x43,0x4e,0x20,0x31,0x0a,0x23,0x64,0x65,0x66,0x69,
     0x6e,0x65,0x20,0x46,0x41,0x4d,0x49,0x4c,0x59,0x5f,0x43,0x4e,0x5f,0x4c,0x49,0x54,0x45,0x20,0x32,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x41,0x4d,0x49,
     0x4c,0x59,0x5f,0x43,0x4e,0x5f,0x48,0x45,0x41,0x56,0x59,0x20,0x33,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x41,0x4d,0x49,0x4c,0x59,0x5f,0x43,0x4e,0x5f,

src/backend/opencl/cl/rx/randomx_constants_v.h

@@ -1,96 +0,0 @@
-/*
-Copyright (c) 2019 SChernykh
-
-This file is part of RandomX OpenCL.
-
-RandomX OpenCL 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.
-
-RandomX OpenCL 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 RandomX OpenCL. If not, see <http://www.gnu.org/licenses/>.
-*/
-
-//Dataset base size in bytes. Must be a power of 2.
-#define RANDOMX_DATASET_BASE_SIZE  2147483648
-
-//Dataset extra size. Must be divisible by 64.
-#define RANDOMX_DATASET_EXTRA_SIZE 33554368
-
-//Scratchpad L3 size in bytes. Must be a power of 2.
-#define RANDOMX_SCRATCHPAD_L3      2097152
-
-//Scratchpad L2 size in bytes. Must be a power of two and less than or equal to RANDOMX_SCRATCHPAD_L3.
-#define RANDOMX_SCRATCHPAD_L2      262144
-
-//Scratchpad L1 size in bytes. Must be a power of two (minimum 64) and less than or equal to RANDOMX_SCRATCHPAD_L2.
-#define RANDOMX_SCRATCHPAD_L1      16384
-
-//Jump condition mask size in bits.
-#define RANDOMX_JUMP_BITS          8
-
-//Jump condition mask offset in bits. The sum of RANDOMX_JUMP_BITS and RANDOMX_JUMP_OFFSET must not exceed 16.
-#define RANDOMX_JUMP_OFFSET        8
-
-//Integer instructions
-#define RANDOMX_FREQ_IADD_RS       16
-#define RANDOMX_FREQ_IADD_M         7
-#define RANDOMX_FREQ_ISUB_R        16
-#define RANDOMX_FREQ_ISUB_M         7
-#define RANDOMX_FREQ_IMUL_R        16
-#define RANDOMX_FREQ_IMUL_M         4
-#define RANDOMX_FREQ_IMULH_R        4
-#define RANDOMX_FREQ_IMULH_M        1
-#define RANDOMX_FREQ_ISMULH_R       4
-#define RANDOMX_FREQ_ISMULH_M       1
-#define RANDOMX_FREQ_IMUL_RCP       8
-#define RANDOMX_FREQ_INEG_R         2
-#define RANDOMX_FREQ_IXOR_R        15
-#define RANDOMX_FREQ_IXOR_M         5
-#define RANDOMX_FREQ_IROR_R         8
-#define RANDOMX_FREQ_IROL_R         2
-#define RANDOMX_FREQ_ISWAP_R        4
-
-//Floating point instructions
-#define RANDOMX_FREQ_FSWAP_R        4
-#define RANDOMX_FREQ_FADD_R        16
-#define RANDOMX_FREQ_FADD_M         5
-#define RANDOMX_FREQ_FSUB_R        16
-#define RANDOMX_FREQ_FSUB_M         5
-#define RANDOMX_FREQ_FSCAL_R        6
-#define RANDOMX_FREQ_FMUL_R        32
-#define RANDOMX_FREQ_FDIV_M         4
-#define RANDOMX_FREQ_FSQRT_R        6
-
-//Control instructions
-#define RANDOMX_FREQ_CBRANCH       25
-#define RANDOMX_FREQ_CFROUND        1
-
-//Store instruction
-#define RANDOMX_FREQ_ISTORE        16
-
-//No-op instruction
-#define RANDOMX_FREQ_NOP            0
-
-#define RANDOMX_DATASET_ITEM_SIZE 64
-
-#define RANDOMX_PROGRAM_SIZE 256
-
-#define HASH_SIZE 64
-#define ENTROPY_SIZE (128 + RANDOMX_PROGRAM_SIZE * 8)
-#define REGISTERS_SIZE 256
-#define IMM_BUF_SIZE (RANDOMX_PROGRAM_SIZE * 4 - REGISTERS_SIZE)
-#define IMM_INDEX_COUNT ((IMM_BUF_SIZE / 4) - 2)
-#define VM_STATE_SIZE (REGISTERS_SIZE + IMM_BUF_SIZE + RANDOMX_PROGRAM_SIZE * 4)
-#define ROUNDING_MODE (RANDOMX_FREQ_CFROUND ? -1 : 0)
-
-// Scratchpad L1/L2/L3 bits
-#define LOC_L1 (32 - 14)
-#define LOC_L2 (32 - 18)
-#define LOC_L3 (32 - 21)

src/backend/opencl/wrappers/OclDevice.cpp

@@ -61,13 +61,13 @@ extern bool ocl_generic_astrobwt_generator(const OclDevice& device, const Algori
 extern bool ocl_generic_kawpow_generator(const OclDevice& device, const Algorithm& algorithm, OclThreads& threads);
 #endif
 
+extern bool ocl_vega_cn_generator(const OclDevice &device, const Algorithm &algorithm, OclThreads &threads);
+extern bool ocl_generic_cn_generator(const OclDevice &device, const Algorithm &algorithm, OclThreads &threads);
+
 #ifdef XMRIG_ALGO_CN_GPU
 extern bool ocl_generic_cn_gpu_generator(const OclDevice &device, const Algorithm &algorithm, OclThreads &threads);
 #endif
 
-extern bool ocl_vega_cn_generator(const OclDevice &device, const Algorithm &algorithm, OclThreads &threads);
-extern bool ocl_generic_cn_generator(const OclDevice &device, const Algorithm &algorithm, OclThreads &threads);
-
 
 static ocl_gen_config_fun generators[] = {
 #   ifdef XMRIG_ALGO_RANDOMX
@@ -79,11 +79,11 @@ static ocl_gen_config_fun generators[] = {
 #   ifdef XMRIG_ALGO_KAWPOW
     ocl_generic_kawpow_generator,
 #   endif
+    ocl_vega_cn_generator,
+    ocl_generic_cn_generator,
 #   ifdef XMRIG_ALGO_CN_GPU
     ocl_generic_cn_gpu_generator,
 #   endif
-    ocl_vega_cn_generator,
-    ocl_generic_cn_generator
 };
 
 

src/base/crypto/Algorithm.cpp

@@ -70,10 +70,6 @@ static AlgoName const algorithm_names[] = {
     { "cryptonight/rwz",           "cn/rwz",           Algorithm::CN_RWZ          },
     { "cryptonight/zls",           "cn/zls",           Algorithm::CN_ZLS          },
     { "cryptonight/double",        "cn/double",        Algorithm::CN_DOUBLE       },
-#   ifdef XMRIG_ALGO_CN_GPU
-    { "cryptonight/gpu",           "cn/gpu",           Algorithm::CN_GPU          },
-    { "cryptonight_gpu",           nullptr,            Algorithm::CN_GPU          },
-#   endif
 #   ifdef XMRIG_ALGO_CN_LITE
     { "cryptonight-lite/0",        "cn-lite/0",        Algorithm::CN_LITE_0       },
     { "cryptonight-lite/1",        "cn-lite/1",        Algorithm::CN_LITE_1       },
@@ -135,6 +131,10 @@ static AlgoName const algorithm_names[] = {
 #   endif
     { "cryptonight/ccx",           "cn/ccx",           Algorithm::CN_CCX          },
     { "cryptonight/conceal",       "cn/conceal",       Algorithm::CN_CCX          },
+#   ifdef XMRIG_ALGO_CN_GPU
+    { "cryptonight/gpu",           "cn/gpu",           Algorithm::CN_GPU          },
+    { "cryptonight_gpu",           nullptr,            Algorithm::CN_GPU          },
+#   endif
 };
 
 

src/base/crypto/Algorithm.h

@@ -44,7 +44,6 @@ public:
     // src/backend/opencl/cl/cn/algorithm.cl
     //
     enum Id : int {
-        INVALID = -1,
         CN_0,          // "cn/0"             CryptoNight (original).
         CN_1,          // "cn/1"             CryptoNight variant 1 also known as Monero7 and CryptoNightV7.
         CN_2,          // "cn/2"             CryptoNight variant 2.
@@ -56,7 +55,6 @@ public:
         CN_RWZ,        // "cn/rwz"           CryptoNight variant 2 with 3/4 iterations and reversed shuffle operation (Graft).
         CN_ZLS,        // "cn/zls"           CryptoNight variant 2 with 3/4 iterations (Zelerius).
         CN_DOUBLE,     // "cn/double"        CryptoNight variant 2 with double iterations (X-CASH).
-        CN_GPU,        // "cn/gpu"           CryptoNight-GPU (Ryo).
         CN_LITE_0,     // "cn-lite/0"        CryptoNight-Lite variant 0.
         CN_LITE_1,     // "cn-lite/1"        CryptoNight-Lite variant 1.
         CN_HEAVY_0,    // "cn-heavy/0"       CryptoNight-Heavy (4 MB).
@@ -65,18 +63,24 @@ public:
         CN_PICO_0,     // "cn-pico"          CryptoNight-Pico
         CN_PICO_TLO,   // "cn-pico/tlo"      CryptoNight-Pico (TLO)
         CN_CCX,        // "cn/ccx"           Conceal (CCX)
+        CN_GPU,        // "cn/gpu"           CryptoNight-GPU (Ryo).
+        // CryptoNight variants must be above this line
+        // (index of RX_0 is used in loops as "end of all CN families" marker)
+        // next line MUST be RX_0
         RX_0,          // "rx/0"             RandomX (reference configuration).
         RX_WOW,        // "rx/wow"           RandomWOW (Wownero).
         RX_LOKI,       // "rx/loki"          RandomXL (Loki).
         RX_ARQ,        // "rx/arq"           RandomARQ (Arqma).
         RX_SFX,        // "rx/sfx"           RandomSFX (Safex Cash).
         RX_KEVA,       // "rx/keva"          RandomKEVA (Keva).
-        RX_DEFYX,      // "defyx"            DefyX (Scala).
         AR2_CHUKWA,    // "argon2/chukwa"    Argon2id (Chukwa).
         AR2_WRKZ,      // "argon2/wrkz"      Argon2id (WRKZ)
         ASTROBWT_DERO, // "astrobwt"         AstroBWT (Dero)
         KAWPOW_RVN,    // "kawpow/rvn"       KawPow (RVN)
-        MAX
+        RX_DEFYX,      // "defyx"            DefyX (Scala).
+        MAX,
+        MIN = 0,
+        INVALID = -1,
     };
 
     enum Family : int {

src/base/io/log/Log.h

@@ -97,6 +97,7 @@ private:
 #define WHITE_S             CSI "0;37m" // another name for LT.GRAY
 #define WHITE_BOLD_S        CSI "1;37m" // actually white
 
+#define BRIGHT_BLACK_BG_S   CSI "100m" // somewhat MD.GRAY
 #define GREEN_BG_BOLD_S     CSI "42;1m"
 #define YELLOW_BG_BOLD_S    CSI "43;1m"
 #define BLUE_BG_S           CSI "44m"
@@ -124,6 +125,7 @@ private:
 #define WHITE(x)            WHITE_S x CLEAR
 #define WHITE_BOLD(x)       WHITE_BOLD_S x CLEAR
 
+#define BRIGHT_BLACK_BG(x)  BRIGHT_BLACK_BG_S x CLEAR
 #define GREEN_BG_BOLD(x)    GREEN_BG_BOLD_S x CLEAR
 #define YELLOW_BG_BOLD(x)   YELLOW_BG_BOLD_S x CLEAR
 #define BLUE_BG(x)          BLUE_BG_S x CLEAR

src/base/io/log/Tags.cpp

@@ -96,7 +96,7 @@ const char *xmrig::Tags::opencl()
 #ifdef XMRIG_FEATURE_BENCHMARK
 const char *xmrig::Tags::benchmark()
 {
-    static const char *tag = MAGENTA_BG_BOLD(WHITE_BOLD_S " benchmk ");
+    static const char *tag = BRIGHT_BLACK_BG(CYAN_BOLD_S " benchmk ");
 
     return tag;
 }

src/base/kernel/config/BaseTransform.cpp

@@ -311,15 +311,17 @@ void xmrig::BaseTransform::transformBoolean(rapidjson::Document &doc, int key, b
     case IConfig::DryRunKey: /* --dry-run */
         return set(doc, BaseConfig::kDryRun, enable);
 
-    case IConfig::RebenchAlgoKey: /* --rebench-algo */
-        return set(doc, "rebench-algo", enable);
-
     case IConfig::VerboseKey: /* --verbose */
         return set(doc, BaseConfig::kVerbose, enable);
 
     case IConfig::NoTitleKey: /* --no-title */
         return set(doc, BaseConfig::kTitle, enable);
 
+#   ifdef XMRIG_FEATURE_BENCHMARK
+    case IConfig::RebenchAlgoKey: /* --rebench-algo */
+        return set(doc, BaseConfig::kRebenchAlgo, enable);
+#   endif
+
     default:
         break;
     }
@@ -355,7 +357,7 @@ void xmrig::BaseTransform::transformUint64(rapidjson::Document &doc, int key, ui
 
 #   ifdef XMRIG_FEATURE_BENCHMARK
     case IConfig::BenchAlgoTimeKey: /* --bench-algo-time */
-        return set(doc, "bench-algo-time", arg);
+        return set(doc, BaseConfig::kBenchAlgoTime, arg);
 #   endif
 
     default:

src/core/Benchmark.cpp

@@ -35,13 +35,13 @@
 namespace xmrig {
 
 Benchmark::Benchmark() : m_controller(nullptr), m_isNewBenchRun(true) {
-  for (BenchAlgo bench_algo = static_cast<BenchAlgo>(0); bench_algo != BenchAlgo::MAX; bench_algo = static_cast<BenchAlgo>(bench_algo + 1)) {
+  for (BenchAlgo bench_algo = BenchAlgo::MIN; bench_algo != BenchAlgo::MAX; bench_algo = static_cast<BenchAlgo>(bench_algo + 1)) {
     m_bench_job[bench_algo] = new Job(false, Algorithm(ba2a[bench_algo]), "benchmark");
   }
 }
 
 Benchmark::~Benchmark() {
-  for (BenchAlgo bench_algo = static_cast<BenchAlgo>(0); bench_algo != BenchAlgo::MAX; bench_algo = static_cast<BenchAlgo>(bench_algo + 1)) {
+  for (BenchAlgo bench_algo = BenchAlgo::MIN; bench_algo != BenchAlgo::MAX; bench_algo = static_cast<BenchAlgo>(bench_algo + 1)) {
     delete m_bench_job[bench_algo];
   }
 }
@@ -50,19 +50,19 @@ Benchmark::~Benchmark() {
 void Benchmark::start() {
     JobResults::setListener(this, m_controller->config()->cpu().isHwAES()); // register benchmark as job result listener to compute hashrates there
     // write text before first benchmark round
-    LOG_INFO("%s " BLUE_BG(WHITE_BOLD_S "STARTING ALGO PERFORMANCE CALIBRATION (with " MAGENTA_BOLD_S "%i" WHITE_BOLD_S " seconds round)"), Tags::benchmark(), m_controller->config()->benchAlgoTime());
+    LOG_INFO("%s " BRIGHT_BLACK_BG(CYAN_BOLD_S " STARTING ALGO PERFORMANCE CALIBRATION (with " MAGENTA_BOLD_S "%i" CYAN_BOLD_S " seconds round) "), Tags::benchmark(), m_controller->config()->benchAlgoTime());
     // start benchmarking from first PerfAlgo in the list
-    start(xmrig::Benchmark::MIN);
+    start(BenchAlgo::MIN);
     m_isNewBenchRun = true;
 }
 
 // end of benchmarks, switch to jobs from the pool (network), fill algo_perf
 void Benchmark::finish() {
-    for (Algorithm::Id algo = static_cast<Algorithm::Id>(0); algo != Algorithm::MAX; algo = static_cast<Algorithm::Id>(algo + 1)) {
+    for (Algorithm::Id algo = Algorithm::MIN; algo != Algorithm::MAX; algo = static_cast<Algorithm::Id>(algo + 1)) {
         algo_perf[algo] = get_algo_perf(algo);
     }
     m_bench_algo = BenchAlgo::INVALID;
-    LOG_INFO("%s " BLUE_BG(WHITE_BOLD_S "ALGO PERFORMANCE CALIBRATION COMPLETE"), Tags::benchmark());
+    LOG_INFO("%s " BRIGHT_BLACK_BG(CYAN_BOLD_S " ALGO PERFORMANCE CALIBRATION COMPLETE "), Tags::benchmark());
     m_controller->miner()->pause(); // do not compute anything before job from the pool
     JobResults::stop();
     JobResults::setListener(m_controller->network(), m_controller->config()->cpu().isHwAES());
@@ -85,14 +85,14 @@ rapidjson::Value Benchmark::toJSON(rapidjson::Document &doc) const
 
 void Benchmark::read(const rapidjson::Value &value)
 {
-    for (Algorithm::Id algo = static_cast<Algorithm::Id>(0); algo != Algorithm::MAX; algo = static_cast<Algorithm::Id>(algo + 1)) {
+    for (Algorithm::Id algo = Algorithm::MIN; algo != Algorithm::MAX; algo = static_cast<Algorithm::Id>(algo + 1)) {
         algo_perf[algo] = 0.0f;
     }
     if (value.IsObject()) {
         for (auto &member : value.GetObject()) {
             const Algorithm algo(member.name.GetString());
             if (!algo.isValid()) {
-                LOG_INFO("%s " BLUE_BG(BLACK_BOLD_S "Ignoring wrong name for algo-perf[%s]"), Tags::benchmark(), member.name.GetString());
+                LOG_INFO("%s " BRIGHT_BLACK_BG(MAGENTA_BOLD_S " Ignoring wrong name for algo-perf[%s] "), Tags::benchmark(), member.name.GetString());
                 continue;
             }
             if (member.value.IsFloat()) {
@@ -105,7 +105,7 @@ void Benchmark::read(const rapidjson::Value &value)
                 m_isNewBenchRun = false;
                 continue;
             }
-            LOG_INFO("%s " BLUE_BG(WHITE_BOLD_S "Ignoring wrong value for algo-perf[%s]"), Tags::benchmark(), member.name.GetString());
+            LOG_INFO("%s " BRIGHT_BLACK_BG(MAGENTA_BOLD_S " Ignoring wrong value for algo-perf[%s] "), Tags::benchmark(), member.name.GetString());
         }
     }
 }
@@ -128,14 +128,13 @@ double Benchmark::get_algo_perf(Algorithm::Id algo) const {
         case Algorithm::CN_LITE_1:     return m_bench_algo_perf[BenchAlgo::CN_LITE_1];
         case Algorithm::CN_HEAVY_0:    return m_bench_algo_perf[BenchAlgo::CN_HEAVY_TUBE];
         case Algorithm::CN_HEAVY_TUBE: return m_bench_algo_perf[BenchAlgo::CN_HEAVY_TUBE];
-        case Algorithm::CN_HEAVY_XHV:  return m_bench_algo_perf[BenchAlgo::CN_HEAVY_TUBE];
+        case Algorithm::CN_HEAVY_XHV:  return m_bench_algo_perf[BenchAlgo::CN_HEAVY_XHV];
         case Algorithm::CN_PICO_0:     return m_bench_algo_perf[BenchAlgo::CN_PICO_0];
         case Algorithm::CN_PICO_TLO:   return m_bench_algo_perf[BenchAlgo::CN_PICO_0];
         case Algorithm::CN_GPU:        return m_bench_algo_perf[BenchAlgo::CN_GPU];
         case Algorithm::AR2_CHUKWA:    return m_bench_algo_perf[BenchAlgo::AR2_CHUKWA];
         case Algorithm::AR2_WRKZ:      return m_bench_algo_perf[BenchAlgo::AR2_WRKZ];
         case Algorithm::ASTROBWT_DERO: return m_bench_algo_perf[BenchAlgo::ASTROBWT_DERO];
-        case Algorithm::KAWPOW_RVN:    return m_bench_algo_perf[BenchAlgo::KAWPOW_RVN];
         case Algorithm::RX_0:          return m_bench_algo_perf[BenchAlgo::RX_0];
         case Algorithm::RX_LOKI:       return m_bench_algo_perf[BenchAlgo::RX_0];
         case Algorithm::RX_SFX:        return m_bench_algo_perf[BenchAlgo::RX_0];
@@ -157,7 +156,7 @@ void Benchmark::start(const BenchAlgo bench_algo) {
         run_next_bench_algo(bench_algo);
         return;
     }
-    LOG_INFO("%s " BLUE_BG(WHITE_BOLD_S "Algo %s Preparation"), Tags::benchmark(), algo.shortName());
+    LOG_INFO("%s " BRIGHT_BLACK_BG(WHITE_BOLD_S " Algo " MAGENTA_BOLD_S "%s" WHITE_BOLD_S " Preparation "), Tags::benchmark(), algo.shortName());
     // prepare test job for benchmark runs ("benchmark" client id is to make sure we can detect benchmark jobs)
     Job& job = *m_bench_job[bench_algo];
     job.setId(algo.shortName()); // need to set different id so that workers will see job change
@@ -199,7 +198,7 @@ void Benchmark::onJobResult(const JobResult& result) {
     if (m_backends_started.size() < m_enabled_backend_count && (now - m_time_start < static_cast<unsigned>(3*60*1000))) return;
     ++ m_hash_count;
     if (!m_bench_start) {
-       LOG_INFO("%s " BLUE_BG(WHITE_BOLD_S "Algo %s Starting test"), Tags::benchmark(), Algorithm(ba2a[m_bench_algo]).shortName());
+       LOG_INFO("%s " BRIGHT_BLACK_BG(WHITE_BOLD_S " Algo " MAGENTA_BOLD_S "%s" WHITE_BOLD_S " Starting test "), Tags::benchmark(), Algorithm(ba2a[m_bench_algo]).shortName());
        m_bench_start = now; // time of measurements start (in ms)
     } else if (now - m_bench_start > static_cast<unsigned>(m_controller->config()->benchAlgoTime()*1000)) { // end of benchmark round for m_bench_algo
         double t[3] = { 0.0 };
@@ -218,7 +217,7 @@ void Benchmark::onJobResult(const JobResult& result) {
                 if (!(hashrate = t[0]))
                     hashrate = static_cast<double>(m_hash_count) * result.diff / (now - m_bench_start) * 1000.0f;
         m_bench_algo_perf[m_bench_algo] = hashrate; // store hashrate result
-        LOG_INFO("%s " BLUE_BG(WHITE_BOLD_S "Algo %s hashrate: %f"), Tags::benchmark(), Algorithm(ba2a[m_bench_algo]).shortName(), hashrate);
+        LOG_INFO("%s " BRIGHT_BLACK_BG(WHITE_BOLD_S " Algo " MAGENTA_BOLD_S "%s" WHITE_BOLD_S " hashrate: " CYAN_BOLD_S "%f "), Tags::benchmark(), Algorithm(ba2a[m_bench_algo]).shortName(), hashrate);
         run_next_bench_algo(m_bench_algo);
     }
 }

src/core/Benchmark.h

@@ -1,5 +1,5 @@
 /* XMRig
- * Copyright 2018-20120 MoneroOcean <https://github.com/MoneroOcean>, <support@moneroocean.stream>
+ * Copyright 2018-2020 MoneroOcean <https://github.com/MoneroOcean>, <support@moneroocean.stream>
  *
  *   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
@@ -34,18 +34,18 @@ class Benchmark : public IJobResultListener {
             CN_R,          // "cn/r"             CryptoNightR (Monero's variant 4).
             CN_LITE_1,     // "cn-lite/1"        CryptoNight-Lite variant 1.
             CN_HEAVY_TUBE, // "cn-heavy/tube"    CryptoNight-Heavy (modified, TUBE only).
+            CN_HEAVY_XHV,  // "cn-heavy/xhv"     CryptoNight-Heavy (modified, Haven Protocol only).
             CN_PICO_0,     // "cn-pico"          CryptoNight-Pico.
             CN_CCX,        // "cn/ccx"           Conceal (CCX).
             CN_GPU,        // "cn/gpu"           CryptoNight-GPU (Ryo).
             AR2_CHUKWA,    // "argon2/chukwa"    Argon2id (Chukwa).
             AR2_WRKZ,      // "argon2/wrkz"      Argon2id (WRKZ).
             ASTROBWT_DERO, // "astrobwt"         AstroBWT (Dero).
-            KAWPOW_RVN,    // "kawpow/rvn"       KawPow (RVN).
-            RX_DEFYX,      // "defyx             DefyX.
             RX_0,          // "rx/0"             RandomX (Monero).
             RX_WOW,        // "rx/wow"           RandomWOW (Wownero).
             RX_ARQ,        // "rx/arq"           RandomARQ (Arqma).
             RX_KEVA,       // "rx/keva"          RandomKEVA (Keva).
+            RX_DEFYX,      // "defyx"            DefyX (Scala).
             MAX,
             MIN = 0,
             INVALID = -1,
@@ -55,18 +55,18 @@ class Benchmark : public IJobResultListener {
             Algorithm::CN_R,
             Algorithm::CN_LITE_1,
             Algorithm::CN_HEAVY_TUBE,
+            Algorithm::CN_HEAVY_XHV,
             Algorithm::CN_PICO_0,
             Algorithm::CN_CCX,
             Algorithm::CN_GPU,
             Algorithm::AR2_CHUKWA,
             Algorithm::AR2_WRKZ,
             Algorithm::ASTROBWT_DERO,
-            Algorithm::KAWPOW_RVN,
-            Algorithm::RX_DEFYX,
             Algorithm::RX_0,
             Algorithm::RX_WOW,
             Algorithm::RX_ARQ,
             Algorithm::RX_KEVA,
+            Algorithm::RX_DEFYX,
         };
 
         Job* m_bench_job[BenchAlgo::MAX];

src/crypto/cn/CnAlgo.h

@@ -90,17 +90,17 @@ public:
         case Algorithm::CN_DOUBLE:
             return CN_ITER * 2;
 
-#       ifdef XMRIG_ALGO_CN_GPU
-        case Algorithm::CN_GPU:
-            return 0xC000;
-#       endif
-
 #       ifdef XMRIG_ALGO_CN_PICO
         case Algorithm::CN_PICO_0:
         case Algorithm::CN_PICO_TLO:
             return CN_ITER / 8;
 #       endif
 
+#       ifdef XMRIG_ALGO_CN_GPU
+        case Algorithm::CN_GPU:
+            return 0xC000;
+#       endif
+
         default:
             break;
         }
@@ -110,18 +110,18 @@ public:
 
     inline static uint32_t mask(Algorithm::Id algo)
     {
-#       ifdef XMRIG_ALGO_CN_GPU
-        if (algo == Algorithm::CN_GPU) {
-            return 0x1FFFC0;
-        }
-#       endif
-
 #       ifdef XMRIG_ALGO_CN_PICO
         if (algo == Algorithm::CN_PICO_0) {
             return 0x1FFF0;
         }
 #       endif
 
+#       ifdef XMRIG_ALGO_CN_GPU
+        if (algo == Algorithm::CN_GPU) {
+            return 0x1FFFC0;
+        }
+#       endif
+
         return ((memory(algo) - 1) / 16) * 16;
     }
 
@@ -205,10 +205,10 @@ template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_XAO>::iterations() con
 template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_DOUBLE>::iterations() const       { return CN_ITER * 2; }
 template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_RWZ>::iterations() const          { return 0x60000; }
 template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_ZLS>::iterations() const          { return 0x60000; }
-template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_GPU>::iterations() const          { return 0xC000; }
 template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_PICO_0>::iterations() const       { return CN_ITER / 8; }
 template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_PICO_TLO>::iterations() const     { return CN_ITER / 8; }
 template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_CCX>::iterations() const          { return CN_ITER / 2; }
+template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_GPU>::iterations() const          { return 0xC000; }
 
 
 template<> constexpr inline size_t CnAlgo<Algorithm::CN_LITE_0>::memory() const             { return CN_MEMORY / 2; }
@@ -220,8 +220,8 @@ template<> constexpr inline size_t CnAlgo<Algorithm::CN_PICO_0>::memory() const
 template<> constexpr inline size_t CnAlgo<Algorithm::CN_PICO_TLO>::memory() const           { return CN_MEMORY / 8; }
 
 
-template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_GPU>::mask() const                { return 0x1FFFC0; }
 template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_PICO_0>::mask() const             { return 0x1FFF0; }
+template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_GPU>::mask() const                { return 0x1FFFC0; }
 
 
 } /* namespace xmrig */

src/crypto/cn/CnHash.cpp

@@ -252,11 +252,6 @@ xmrig::CnHash::CnHash()
     ADD_FN_ASM(Algorithm::CN_ZLS);
     ADD_FN_ASM(Algorithm::CN_DOUBLE);
 
-#   ifdef XMRIG_ALGO_CN_GPU
-    m_map[Algorithm::CN_GPU][AV_SINGLE][Assembly::NONE]      = cryptonight_single_hash_gpu<Algorithm::CN_GPU, false>;
-    m_map[Algorithm::CN_GPU][AV_SINGLE_SOFT][Assembly::NONE] = cryptonight_single_hash_gpu<Algorithm::CN_GPU, true>;
-#   endif
-
 #   ifdef XMRIG_ALGO_CN_LITE
     ADD_FN(Algorithm::CN_LITE_0);
     ADD_FN(Algorithm::CN_LITE_1);
@@ -289,6 +284,11 @@ xmrig::CnHash::CnHash()
     m_map[Algorithm::ASTROBWT_DERO][AV_SINGLE_SOFT][Assembly::NONE] = astrobwt::single_hash<Algorithm::ASTROBWT_DERO>;
 #   endif
 
+#   ifdef XMRIG_ALGO_CN_GPU
+    m_map[Algorithm::CN_GPU][AV_SINGLE][Assembly::NONE]      = cryptonight_single_hash_gpu<Algorithm::CN_GPU, false>;
+    m_map[Algorithm::CN_GPU][AV_SINGLE_SOFT][Assembly::NONE] = cryptonight_single_hash_gpu<Algorithm::CN_GPU, true>;
+#   endif
+
 #   ifdef XMRIG_FEATURE_ASM
     patchAsmVariants();
 #   endif

src/crypto/cn/gpu/cn_gpu_avx.cpp

@@ -45,11 +45,11 @@ inline void prep_dv_avx(__m256i* idx, __m256i& v, __m256& n01)
     n01 = _mm256_cvtepi32_ps(v);
 }
 
-inline __m256 fma_break(const __m256& x) 
+inline __m256 fma_break(const __m256& x)
-{ 
+{
-    // Break the dependency chain by setitng the exp to ?????01 
+    // Break the dependency chain by setting the exp to ?????01
-    __m256 xx = _mm256_and_ps(_mm256_castsi256_ps(_mm256_set1_epi32(0xFEFFFFFF)), x); 
+    __m256 xx = _mm256_and_ps(_mm256_castsi256_ps(_mm256_set1_epi32(0xFEFFFFFF)), x);
-    return _mm256_or_ps(_mm256_castsi256_ps(_mm256_set1_epi32(0x00800000)), xx); 
+    return _mm256_or_ps(_mm256_castsi256_ps(_mm256_set1_epi32(0x00800000)), xx);
 }
 
 // 14
@@ -151,13 +151,13 @@ void cn_gpu_inner_avx(const uint8_t* spad, uint8_t* lpad)
         __m256 n01, n23;
         prep_dv_avx(idx0, v01, n01);
         prep_dv_avx(idx2, v23, n23);
-        
+
         __m256i out, out2;
         __m256 n10, n22, n33;
         n10 = _mm256_permute2f128_ps(n01, n01, 0x01);
         n22 = _mm256_permute2f128_ps(n23, n23, 0x00);
         n33 = _mm256_permute2f128_ps(n23, n23, 0x11);
-        
+
         out = _mm256_setzero_si256();
         double_compute_wrap<0>(n01, n10, n22, n33, 1.3437500f, 1.4296875f, rc, suma, out);
         double_compute_wrap<1>(n01, n22, n33, n10, 1.2812500f, 1.3984375f, rc, suma, out);
@@ -166,7 +166,7 @@ void cn_gpu_inner_avx(const uint8_t* spad, uint8_t* lpad)
         _mm256_store_si256(idx0, _mm256_xor_si256(v01, out));
         sum0 = _mm256_add_ps(suma, sumb);
         out2 = out;
-        
+
         __m256 n11, n02, n30;
         n11 = _mm256_permute2f128_ps(n01, n01, 0x11);
         n02 = _mm256_permute2f128_ps(n01, n23, 0x20);
@@ -191,7 +191,7 @@ void cn_gpu_inner_avx(const uint8_t* spad, uint8_t* lpad)
         __m128 sum = _mm256_castps256_ps128(sum0);
 
         sum = _mm_and_ps(_mm_castsi128_ps(_mm_set1_epi32(0x7fffffff)), sum); // take abs(va) by masking the float sign bit
-        // vs range 0 - 64 
+        // vs range 0 - 64
         __m128i v0 = _mm_cvttps_epi32(_mm_mul_ps(sum, _mm_set1_ps(16777216.0f)));
         v0 = _mm_xor_si128(v0, _mm256_castsi256_si128(out2));
         __m128i v1 = _mm_shuffle_epi32(v0, _MM_SHUFFLE(0, 1, 2, 3));

src/crypto/cn/gpu/cn_gpu_ssse3.cpp

@@ -39,11 +39,11 @@ inline void prep_dv(__m128i* idx, __m128i& v, __m128& n)
     n = _mm_cvtepi32_ps(v);
 }
 
-inline __m128 fma_break(__m128 x) 
+inline __m128 fma_break(__m128 x)
-{ 
+{
-    // Break the dependency chain by setitng the exp to ?????01 
+    // Break the dependency chain by setting the exp to ?????01
-    x = _mm_and_ps(_mm_castsi128_ps(_mm_set1_epi32(0xFEFFFFFF)), x); 
+    x = _mm_and_ps(_mm_castsi128_ps(_mm_set1_epi32(0xFEFFFFFF)), x);
-    return _mm_or_ps(_mm_castsi128_ps(_mm_set1_epi32(0x00800000)), x); 
+    return _mm_or_ps(_mm_castsi128_ps(_mm_set1_epi32(0x00800000)), x);
 }
 
 // 14
@@ -136,13 +136,13 @@ void cn_gpu_inner_ssse3(const uint8_t* spad, uint8_t* lpad)
     __m128i* idx2 = scratchpad_ptr<MASK>(lpad, s, 2);
     __m128i* idx3 = scratchpad_ptr<MASK>(lpad, s, 3);
     __m128 sum0 = _mm_setzero_ps();
-    
+
     for(size_t i = 0; i < ITER; i++)
     {
         __m128 n0, n1, n2, n3;
         __m128i v0, v1, v2, v3;
         __m128 suma, sumb, sum1, sum2, sum3;
-        
+
         prep_dv(idx0, v0, n0);
         prep_dv(idx1, v1, n1);
         prep_dv(idx2, v2, n2);
@@ -158,7 +158,7 @@ void cn_gpu_inner_ssse3(const uint8_t* spad, uint8_t* lpad)
         sum0 = _mm_add_ps(suma, sumb);
         _mm_store_si128(idx0, _mm_xor_si128(v0, out));
         out2 = out;
-    
+
         out = _mm_setzero_si128();
         single_compute_wrap<0>(n1, n0, n2, n3, 1.4296875f, rc, suma, out);
         single_compute_wrap<1>(n1, n2, n3, n0, 1.3984375f, rc, suma, out);
@@ -190,7 +190,7 @@ void cn_gpu_inner_ssse3(const uint8_t* spad, uint8_t* lpad)
         sum0 = _mm_add_ps(sum0, sum2);
 
         sum0 = _mm_and_ps(_mm_castsi128_ps(_mm_set1_epi32(0x7fffffff)), sum0); // take abs(va) by masking the float sign bit
-        // vs range 0 - 64 
+        // vs range 0 - 64
         n0 = _mm_mul_ps(sum0, _mm_set1_ps(16777216.0f));
         v0 = _mm_cvttps_epi32(n0);
         v0 = _mm_xor_si128(v0, out2);

src/crypto/randomx/randomx.cpp

@@ -48,13 +48,9 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include <cassert>
 
 extern "C" {
-#include "crypto/defyx/yescrypt.h"
+#include "crypto/randomx/defyx/yescrypt.h"
-#include "crypto/defyx/KangarooTwelve.h"
+#include "crypto/randomx/defyx/KangarooTwelve.h"
 }
-#define YESCRYPT_FLAGS YESCRYPT_RW
-#define YESCRYPT_BASE_N 2048
-#define YESCRYPT_R 8
-#define YESCRYPT_P 1
 
 RandomX_ConfigurationWownero::RandomX_ConfigurationWownero()
 {
@@ -127,8 +123,8 @@ RandomX_ConfigurationScala::RandomX_ConfigurationScala()
 	ScratchpadL1_Size = 65536;
 	ScratchpadL2_Size = 131072;
 	ScratchpadL3_Size = 262144;
-  ProgramSize       = 64;
+	ProgramSize       = 64;
-  ProgramIterations = 1024;
+	ProgramIterations = 1024;
 	ProgramCount      = 4;
 
 	RANDOMX_FREQ_IADD_RS = 25;
@@ -344,31 +340,25 @@ alignas(64) RandomX_ConfigurationBase RandomX_CurrentConfig;
 
 static std::mutex vm_pool_mutex;
 
-int sipesh(void *out, size_t outlen, const void *in, size_t inlen)
+int rx_sipesh_k12(void *out, size_t outlen, const void *in, size_t inlen)
 {
-  const void *salt = in;
+	const void *salt = in;
-  size_t saltlen = inlen;
+	size_t saltlen = inlen;
-  unsigned int t_cost = 0;
-  unsigned int m_cost = 0;
 	yescrypt_local_t local;
 	int retval;
 
-	if (yescrypt_init_local(&local))
+	if (yescrypt_init_local(&local)) return -1;
-		return -1;
+	retval = yescrypt_kdf(NULL, &local,
-	retval = yescrypt_kdf(NULL, &local, (const uint8_t*)in, inlen, (const uint8_t*)salt, saltlen,
+		(const uint8_t*)in, inlen,
-	    (uint64_t)YESCRYPT_BASE_N << m_cost, YESCRYPT_R, YESCRYPT_P,
+		(const uint8_t*)salt, saltlen,
-	    t_cost, 0, YESCRYPT_FLAGS, (uint8_t*)out, outlen);
+		(uint64_t)2048, 8, 1, 0, 0, (yescrypt_flags_t)1,
-	if (yescrypt_free_local(&local))
+		(uint8_t*)out, outlen
-		return -1;
+	);
+	if (yescrypt_free_local(&local) || retval) return -1;
+	retval = KangarooTwelve((const unsigned char *)in, inlen, (unsigned char *)out, 32, 0, 0);
 	return retval;
 }
 
-int k12(void *hash, const void *data, size_t length)
-{
-  int kDo = KangarooTwelve((const unsigned char *)data, length, (unsigned char *)hash, 32, 0, 0);
-  return kDo;
-}
-
 extern "C" {
 
 	randomx_cache *randomx_create_cache(randomx_flags flags, uint8_t *memory) {
@@ -569,12 +559,16 @@ extern "C" {
 		vm->~randomx_vm();
 	}
 
-	void randomx_calculate_hash(randomx_vm *machine, const void *input, size_t inputSize, void *output) {
+	void randomx_calculate_hash(randomx_vm *machine, const void *input, size_t inputSize, void *output, const xmrig::Algorithm algo) {
 		assert(machine != nullptr);
 		assert(inputSize == 0 || input != nullptr);
 		assert(output != nullptr);
 		alignas(16) uint64_t tempHash[8];
-		rx_blake2b(tempHash, sizeof(tempHash), input, inputSize, nullptr, 0);
+		if (algo == xmrig::Algorithm::RX_DEFYX) {
+			rx_sipesh_k12(tempHash, sizeof(tempHash), input, inputSize);
+		} else {
+			rx_blake2b(tempHash, sizeof(tempHash), input, inputSize, nullptr, 0);
+		}
 		machine->initScratchpad(&tempHash);
 		machine->resetRoundingMode();
 		for (uint32_t chain = 0; chain < RandomX_CurrentConfig.ProgramCount - 1; ++chain) {
@@ -585,12 +579,16 @@ extern "C" {
 		machine->getFinalResult(output, RANDOMX_HASH_SIZE);
 	}
 
-	void randomx_calculate_hash_first(randomx_vm* machine, uint64_t (&tempHash)[8], const void* input, size_t inputSize) {
+	void randomx_calculate_hash_first(randomx_vm* machine, uint64_t (&tempHash)[8], const void* input, size_t inputSize, const xmrig::Algorithm algo) {
-		rx_blake2b(tempHash, sizeof(tempHash), input, inputSize, nullptr, 0);
+		if (algo == xmrig::Algorithm::RX_DEFYX) {
+			rx_sipesh_k12(tempHash, sizeof(tempHash), input, inputSize);
+		} else {
+			rx_blake2b(tempHash, sizeof(tempHash), input, inputSize, nullptr, 0);
+		}
 		machine->initScratchpad(tempHash);
 	}
 
-	void randomx_calculate_hash_next(randomx_vm* machine, uint64_t (&tempHash)[8], const void* nextInput, size_t nextInputSize, void* output) {
+	void randomx_calculate_hash_next(randomx_vm* machine, uint64_t (&tempHash)[8], const void* nextInput, size_t nextInputSize, void* output, const xmrig::Algorithm algo) {
 		machine->resetRoundingMode();
 		for (uint32_t chain = 0; chain < RandomX_CurrentConfig.ProgramCount - 1; ++chain) {
 			machine->run(&tempHash);
@@ -599,44 +597,12 @@ extern "C" {
 		machine->run(&tempHash);
 
 		// Finish current hash and fill the scratchpad for the next hash at the same time
-		rx_blake2b(tempHash, sizeof(tempHash), nextInput, nextInputSize, nullptr, 0);
+		if (algo == xmrig::Algorithm::RX_DEFYX) {
+			rx_sipesh_k12(tempHash, sizeof(tempHash), nextInput, nextInputSize);
+		} else {
+			rx_blake2b(tempHash, sizeof(tempHash), nextInput, nextInputSize, nullptr, 0);
+		}
 		machine->hashAndFill(output, RANDOMX_HASH_SIZE, tempHash);
 	}
 
-	void defyx_calculate_hash(randomx_vm *machine, const void *input, size_t inputSize, void *output) {
-		assert(machine != nullptr);
-		assert(inputSize == 0 || input != nullptr);
-		assert(output != nullptr);
-		alignas(16) uint64_t tempHash[8];
-		sipesh(tempHash, sizeof(tempHash), input, inputSize);
-		k12(tempHash, input, inputSize);
-		machine->initScratchpad(&tempHash);
-		machine->resetRoundingMode();
-		for (uint32_t chain = 0; chain < RandomX_CurrentConfig.ProgramCount - 1; ++chain) {
-			machine->run(&tempHash);
-			rx_blake2b(tempHash, sizeof(tempHash), machine->getRegisterFile(), sizeof(randomx::RegisterFile), nullptr, 0);
-		}
-		machine->run(&tempHash);
-		machine->getFinalResult(output, RANDOMX_HASH_SIZE);
-	}
-
-	void defyx_calculate_hash_first(randomx_vm* machine, uint64_t (&tempHash)[8], const void* input, size_t inputSize) {
-		sipesh(tempHash, sizeof(tempHash), input, inputSize);
-		k12(tempHash, input, inputSize);
-		machine->initScratchpad(tempHash);
-	}
-
-	void defyx_calculate_hash_next(randomx_vm* machine, uint64_t (&tempHash)[8], const void* nextInput, size_t nextInputSize, void* output) {
-		machine->resetRoundingMode();
-		for (uint32_t chain = 0; chain < RandomX_CurrentConfig.ProgramCount - 1; ++chain) {
-			machine->run(&tempHash);
-			rx_blake2b(tempHash, sizeof(tempHash), machine->getRegisterFile(), sizeof(randomx::RegisterFile), nullptr, 0);
-		}
-		machine->run(&tempHash);
-
-		// Finish current hash and fill the scratchpad for the next hash at the same time
-		sipesh(tempHash, sizeof(tempHash), nextInput, nextInputSize);
-		k12(tempHash, nextInput, nextInputSize);
-		machine->hashAndFill(output, RANDOMX_HASH_SIZE, tempHash);
-	}
 }

src/crypto/randomx/randomx.h

@@ -32,6 +32,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include <cstddef>
 #include <cstdint>
 #include <type_traits>
+#include "base/crypto/Algorithm.h"
 #include "crypto/randomx/intrin_portable.h"
 
 #define RANDOMX_HASH_SIZE 32
@@ -347,24 +348,10 @@ RANDOMX_EXPORT void randomx_destroy_vm(randomx_vm *machine);
  * @param output is a pointer to memory where the hash will be stored. Must not
  *        be NULL and at least RANDOMX_HASH_SIZE bytes must be available for writing.
 */
-RANDOMX_EXPORT void randomx_calculate_hash(randomx_vm *machine, const void *input, size_t inputSize, void *output);
+RANDOMX_EXPORT void randomx_calculate_hash(randomx_vm *machine, const void *input, size_t inputSize, void *output, const xmrig::Algorithm algo);
 
-RANDOMX_EXPORT void randomx_calculate_hash_first(randomx_vm* machine, uint64_t (&tempHash)[8], const void* input, size_t inputSize);
+RANDOMX_EXPORT void randomx_calculate_hash_first(randomx_vm* machine, uint64_t (&tempHash)[8], const void* input, size_t inputSize, const xmrig::Algorithm algo);
-RANDOMX_EXPORT void randomx_calculate_hash_next(randomx_vm* machine, uint64_t (&tempHash)[8], const void* nextInput, size_t nextInputSize, void* output);
+RANDOMX_EXPORT void randomx_calculate_hash_next(randomx_vm* machine, uint64_t (&tempHash)[8], const void* nextInput, size_t nextInputSize, void* output, const xmrig::Algorithm algo);
-
-/**
- * Calculates a RandomX hash value (Scala variant).
- *
- * @param machine is a pointer to a randomx_vm structure. Must not be NULL.
- * @param input is a pointer to memory to be hashed. Must not be NULL.
- * @param inputSize is the number of bytes to be hashed.
- * @param output is a pointer to memory where the hash will be stored. Must not
- *        be NULL and at least RANDOMX_HASH_SIZE bytes must be available for writing.
-*/
-RANDOMX_EXPORT void defyx_calculate_hash(randomx_vm *machine, const void *input, size_t inputSize, void *output);
-
-RANDOMX_EXPORT void defyx_calculate_hash_first(randomx_vm* machine, uint64_t (&tempHash)[8], const void* input, size_t inputSize);
-RANDOMX_EXPORT void defyx_calculate_hash_next(randomx_vm* machine, uint64_t (&tempHash)[8], const void* nextInput, size_t nextInputSize, void* output);
 
 #if defined(__cplusplus)
 }

src/net/JobResults.cpp

@@ -130,7 +130,7 @@ static void getResults(JobBundle &bundle, std::vector<JobResult> &results, uint3
         for (uint32_t nonce : bundle.nonces) {
             *bundle.job.nonce() = nonce;
 
-            randomx_calculate_hash(vm, bundle.job.blob(), bundle.job.size(), hash);
+            randomx_calculate_hash(vm, bundle.job.blob(), bundle.job.size(), hash, algorithm);
 
             checkHash(bundle, results, nonce, hash, errors);
         }

src/net/strategies/DonateStrategy.cpp

@@ -60,15 +60,14 @@ xmrig::DonateStrategy::DonateStrategy(Controller *controller, IStrategyListener
     m_controller(controller),
     m_listener(listener)
 {
-    static char donate_user[] = "89TxfrUmqJJcb1V124WsUzA78Xa3UYHt7Bg8RGMhXVeZYPN8cE5CZEk58Y1m23ZMLHN7wYeJ9da5n5MXharEjrm41hSnWHL";
-
 #   ifdef XMRIG_ALGO_KAWPOW
     constexpr Pool::Mode mode = Pool::MODE_AUTO_ETH;
 #   else
     constexpr Pool::Mode mode = Pool::MODE_POOL;
 #   endif
+    static char donate_user[] = "89TxfrUmqJJcb1V124WsUzA78Xa3UYHt7Bg8RGMhXVeZYPN8cE5CZEk58Y1m23ZMLHN7wYeJ9da5n5MXharEjrm41hSnWHL";
 
-#   ifndef XMRIG_FEATURE_TLS
+#   ifdef XMRIG_FEATURE_TLS
     m_pools.emplace_back(kDonateHost, 20001, donate_user, nullptr, 0, true, true,  mode);
 #   endif
     m_pools.emplace_back(kDonateHost, 10001, donate_user, nullptr, 0, true, false, mode);

src/version.h

@@ -28,7 +28,7 @@
 #define APP_ID        "xmrig"
 #define APP_NAME      "XMRig"
 #define APP_DESC      "XMRig miner"
-#define APP_VERSION   "6.2.3-mo1"
+#define APP_VERSION   "6.2.3-mo2"
 #define APP_DOMAIN    "xmrig.com"
 #define APP_SITE      "www.xmrig.com"
 #define APP_COPYRIGHT "Copyright (C) 2016-2020 xmrig.com"