v6.2.3

RandomX: tweaked Ryzen code

CHANGELOG.md

@@ -1,3 +1,13 @@
+# v6.2.3
+- [#1745](https://github.com/xmrig/xmrig/pull/1745) AstroBWT: fixed OpenCL compilation on some systems.
+- [#1749](https://github.com/xmrig/xmrig/pull/1749) KawPow: optimized CPU share verification.
+- [#1752](https://github.com/xmrig/xmrig/pull/1752) RandomX: added error message when MSR mod fails.
+- [#1754](https://github.com/xmrig/xmrig/issues/1754) Fixed GPU health readings for pre Vega GPUs on Linux.
+- [#1756](https://github.com/xmrig/xmrig/issues/1756) Added results and connection reports.
+- [#1759](https://github.com/xmrig/xmrig/pull/1759) KawPow: fixed DAG initialization on slower AMD GPUs.
+- [#1763](https://github.com/xmrig/xmrig/pull/1763) KawPow: fixed rare duplicate share errors.
+- [#1766](https://github.com/xmrig/xmrig/pull/1766) RandomX: small speedup on Ryzen CPUs.
+
 # v6.2.2
 - [#1742](https://github.com/xmrig/xmrig/issues/1742) Fixed crash when use HTTP API.
 

src/3rdparty/libethash/ethash_internal.c

@@ -33,6 +33,18 @@
 #include "data_sizes.h"
 #include "base/crypto/sha3.h"
 
+#if defined(_M_X64) || defined(__x86_64__) || defined(__SSE2__)
+	#ifdef __GNUC__
+		#include <x86intrin.h>
+	#else
+		#include <intrin.h>
+	#endif
+
+	#define kp_prefetch(x) _mm_prefetch((x), _MM_HINT_T0);
+#else
+	#define kp_prefetch(x)
+#endif
+
 #define SHA3_256(a, b, c) sha3_HashBuffer(256, SHA3_FLAGS_KECCAK, b, c, a, 32)
 #define SHA3_512(a, b, c) sha3_HashBuffer(512, SHA3_FLAGS_KECCAK, b, c, a, 64)
 
@@ -157,46 +169,53 @@ void ethash_calculate_dag_item_opt(
 	memcpy(ret, init, sizeof(node));
 	ret->words[0] ^= node_index;
 	SHA3_512(ret->bytes, ret->bytes, sizeof(node));
-#if defined(_M_X64) && ENABLE_SSE
-	__m128i const fnv_prime = _mm_set1_epi32(FNV_PRIME);
-	__m128i xmm0 = ret->xmm[0];
-	__m128i xmm1 = ret->xmm[1];
-	__m128i xmm2 = ret->xmm[2];
-	__m128i xmm3 = ret->xmm[3];
-#endif
 
 	for (uint32_t i = 0; i != num_parents; ++i) {
 		uint32_t parent_index = fast_mod(fnv_hash(node_index ^ i, ret->words[i % NODE_WORDS]), light->num_parent_nodes, light->reciprocal, light->increment, light->shift);
 		node const* parent = &cache_nodes[parent_index];
-
-#if defined(_M_X64) && ENABLE_SSE
-		{
-			xmm0 = _mm_mullo_epi32(xmm0, fnv_prime);
-			xmm1 = _mm_mullo_epi32(xmm1, fnv_prime);
-			xmm2 = _mm_mullo_epi32(xmm2, fnv_prime);
-			xmm3 = _mm_mullo_epi32(xmm3, fnv_prime);
-			xmm0 = _mm_xor_si128(xmm0, parent->xmm[0]);
-			xmm1 = _mm_xor_si128(xmm1, parent->xmm[1]);
-			xmm2 = _mm_xor_si128(xmm2, parent->xmm[2]);
-			xmm3 = _mm_xor_si128(xmm3, parent->xmm[3]);
-
-			// have to write to ret as values are used to compute index
-			ret->xmm[0] = xmm0;
-			ret->xmm[1] = xmm1;
-			ret->xmm[2] = xmm2;
-			ret->xmm[3] = xmm3;
+		for (unsigned w = 0; w != NODE_WORDS; ++w) {
+			ret->words[w] = fnv_hash(ret->words[w], parent->words[w]);
 		}
-#else
-		{
-			for (unsigned w = 0; w != NODE_WORDS; ++w) {
-				ret->words[w] = fnv_hash(ret->words[w], parent->words[w]);
-			}
-		}
-#endif
 	}
 	SHA3_512(ret->bytes, ret->bytes, sizeof(node));
 }
 
+void ethash_calculate_dag_item4_opt(
+	node* ret,
+	uint32_t node_index,
+	uint32_t num_parents,
+	ethash_light_t const light
+)
+{
+	node const* cache_nodes = (node const*)light->cache;
+
+	for (size_t i = 0; i < 4; ++i) {
+		node const* init = &cache_nodes[fast_mod(node_index + i, light->num_parent_nodes, light->reciprocal, light->increment, light->shift)];
+		memcpy(ret + i, init, sizeof(node));
+		ret[i].words[0] ^= node_index + i;
+		SHA3_512(ret[i].bytes, ret[i].bytes, sizeof(node));
+	}
+
+	for (uint32_t i = 0; i != num_parents; ++i) {
+		node const* parent[4];
+
+		for (uint32_t j = 0; j < 4; ++j) {
+			const uint32_t parent_index = fast_mod(fnv_hash((node_index + j) ^ i, ret[j].words[i % NODE_WORDS]), light->num_parent_nodes, light->reciprocal, light->increment, light->shift);
+			parent[j] = &cache_nodes[parent_index];
+			kp_prefetch(parent[j]);
+		}
+
+		for (unsigned w = 0; w != NODE_WORDS; ++w) ret[0].words[w] = fnv_hash(ret[0].words[w], parent[0]->words[w]);
+		for (unsigned w = 0; w != NODE_WORDS; ++w) ret[1].words[w] = fnv_hash(ret[1].words[w], parent[1]->words[w]);
+		for (unsigned w = 0; w != NODE_WORDS; ++w) ret[2].words[w] = fnv_hash(ret[2].words[w], parent[2]->words[w]);
+		for (unsigned w = 0; w != NODE_WORDS; ++w) ret[3].words[w] = fnv_hash(ret[3].words[w], parent[3]->words[w]);
+	}
+
+	for (size_t i = 0; i < 4; ++i) {
+		SHA3_512(ret[i].bytes, ret[i].bytes, sizeof(node));
+	}
+}
+
 bool ethash_compute_full_data(
 	void* mem,
 	uint64_t full_size,

src/3rdparty/libethash/ethash_internal.h

@@ -161,6 +161,13 @@ void ethash_calculate_dag_item_opt(
 	ethash_light_t const cache
 );
 
+void ethash_calculate_dag_item4_opt(
+	node* ret,
+	uint32_t node_index,
+	uint32_t num_parents,
+	ethash_light_t const cache
+);
+
 void ethash_quick_hash(
 	ethash_h256_t* return_hash,
 	ethash_h256_t const* header_hash,

src/App.cpp

@@ -37,8 +37,6 @@
 #include "base/kernel/Platform.h"
 #include "core/config/Config.h"
 #include "core/Controller.h"
-#include "core/Miner.h"
-#include "net/Network.h"
 #include "Summary.h"
 #include "version.h"
 
@@ -107,7 +105,7 @@ void xmrig::App::onConsoleCommand(char command)
         close();
     }
     else {
-        m_controller->miner()->execCommand(command);
+        m_controller->execCommand(command);
     }
 }
 

src/Summary.cpp

@@ -161,11 +161,14 @@ static void print_commands(Config *)
 {
     if (Log::isColors()) {
         Log::print(GREEN_BOLD(" * ") WHITE_BOLD("COMMANDS     ") MAGENTA_BG(WHITE_BOLD_S "h") WHITE_BOLD("ashrate, ")
-                                                                     MAGENTA_BG(WHITE_BOLD_S "p") WHITE_BOLD("ause, ")
-                                                                     MAGENTA_BG(WHITE_BOLD_S "r") WHITE_BOLD("esume"));
+                                                                 MAGENTA_BG(WHITE_BOLD_S "p") WHITE_BOLD("ause, ")
+                                                                 MAGENTA_BG(WHITE_BOLD_S "r") WHITE_BOLD("esume, ")
+                                                                 WHITE_BOLD("re") MAGENTA_BG(WHITE_BOLD_S "s") WHITE_BOLD("ults, ")
+                                                                 MAGENTA_BG(WHITE_BOLD_S "c") WHITE_BOLD("onnection")
+                   );
     }
     else {
-        Log::print(" * COMMANDS     'h' hashrate, 'p' pause, 'r' resume");
+        Log::print(" * COMMANDS     'h' hashrate, 'p' pause, 'r' resume, 's' results, 'c' connection");
     }
 }
 

src/backend/common/WorkerJob.h

@@ -139,11 +139,24 @@ inline bool xmrig::WorkerJob<1>::nextRound(uint32_t rounds, uint32_t roundSize)
     }
 
     // Increment higher 32 bits of a 64-bit nonce when lower 32 bits overflow
-    if (!currentJob().isNicehash() && (nonceSize() == sizeof(uint64_t)) && (*n < prev_nonce)) {
-        ++n[1];
+    if (!currentJob().isNicehash() && (nonceSize() == sizeof(uint64_t))) {
+        const bool wrapped = (*n < prev_nonce);
+        const bool wraps_this_round = (static_cast<uint64_t>(*n) + roundSize > (1ULL << 32));
 
-        Job& job = m_jobs[index()];
-        memcpy(job.blob(), blob(), job.size());
+        // Account for the case when starting nonce hasn't wrapped yet, but some nonces in the current round will wrap
+        if (wrapped || wraps_this_round) {
+            // Set lower 32 bits to 0 when higher 32 bits change
+            Nonce::reset(index());
+
+            // Sets *n to 0 and Nonce::m_nonce[index] to the correct next value
+            *n = 0;
+            Nonce::next(index(), *n, rounds * roundSize, currentJob().isNicehash(), &ok);
+
+            ++n[1];
+
+            Job& job = m_jobs[index()];
+            memcpy(job.blob(), blob(), job.size());
+        }
     }
 
     return ok;

src/backend/opencl/cl/astrobwt/BWT.cl

@@ -163,7 +163,7 @@ __kernel void filter(uint32_t nonce, uint32_t bwt_max_size, __global const uint3
 
 		filtered_hashes[index] = nonce + global_id;
 
-		#pragma unroll(8)
+		#pragma unroll 8
 		for (uint32_t i = 0; i < 8; ++i)
 			filtered_hashes[index + i + 1] = hash[i];
 	}
@@ -183,7 +183,7 @@ __kernel void prepare_batch2(__global uint32_t* hashes, __global uint32_t* filte
 	const uint32_t stage2_size = STAGE1_SIZE + (filtered_hash[1] & 0xfffff);
 	data_sizes[global_id] = stage2_size;
 
-	#pragma unroll(8)
+	#pragma unroll 8
 	for (uint32_t i = 0; i < 8; ++i)
 		hash[i] = filtered_hash[i + 1];
 }

src/backend/opencl/cl/astrobwt/salsa20.cl

@@ -82,7 +82,7 @@ __kernel void Salsa20_XORKeyStream(__global const uint32_t* keys, __global uint3
 		uint32_t x14 = j14;
 		uint32_t x15 = j15;
 
-		#pragma unroll(5)
+		#pragma unroll 5
 		for (uint32_t j = 0; j < 10; ++j)
 		{
 			x4  = XOR( x4,ROTATE(PLUS( x0,x12), 7));

src/backend/opencl/cl/astrobwt/sha3.cl

@@ -111,7 +111,7 @@ __kernel void sha3(__global const uint8_t* inputs, __global const uint32_t* inpu
 		++wordIndex;
 		if (wordIndex == 17)
 		{
-			#pragma unroll(ROUNDS)
+			#pragma unroll ROUNDS
 			for (int i = 0; i < ROUNDS; ++i)
 			{
 				C[t] = A[s] ^ A[s+5] ^ A[s+10] ^ A[s+15] ^ A[s+20];
@@ -133,7 +133,7 @@ __kernel void sha3(__global const uint8_t* inputs, __global const uint32_t* inpu
 	A[wordIndex] ^= tail ^ ((uint64_t)(((uint64_t)(0x02 | (1 << 2))) << (tail_size * 8)));
 	A[16] ^= 0x8000000000000000UL;
 
-	#pragma unroll(1)
+	#pragma unroll 1
 	for (int i = 0; i < ROUNDS; ++i)
 	{
 		C[t] = A[s] ^ A[s+5] ^ A[s+10] ^ A[s+15] ^ A[s+20];
@@ -180,7 +180,7 @@ __kernel void sha3_initial(__global const uint8_t* input_data, uint32_t input_si
 	A[wordIndex] ^= (uint64_t)(((uint64_t)(0x02 | (1 << 2))) << (tail_size * 8));
 	A[16] ^= 0x8000000000000000UL;
 
-	#pragma unroll(ROUNDS)
+	#pragma unroll ROUNDS
 	for (int i = 0; i < ROUNDS; ++i)
 	{
 		C[t] = A[s] ^ A[s+5] ^ A[s+10] ^ A[s+15] ^ A[s+20];

src/backend/opencl/runners/OclKawPowRunner.cpp

@@ -157,7 +157,7 @@ void OclKawPowRunner::set(const Job &job, uint8_t *blob)
         const uint32_t dag_words = dag_size / sizeof(node);
         m_calculateDagKernel->setArgs(0, m_lightCache, m_dag, dag_words, m_lightCacheSize / sizeof(node));
 
-        constexpr uint32_t N = 1 << 20;
+        constexpr uint32_t N = 1 << 18;
 
         for (uint32_t start = 0; start < dag_words; start += N) {
             m_calculateDagKernel->setArg(0, sizeof(start), &start);

src/backend/opencl/wrappers/AdlLib_linux.cpp

@@ -1,6 +1,7 @@
 /* XMRig
  * Copyright 2008-2018 Advanced Micro Devices, Inc.
  * Copyright 2018-2020 SChernykh                    <https://github.com/SChernykh>
+ * Copyright 2020      Patrick Bollinger            <https://github.com/pjbollinger>
  * Copyright 2016-2020 XMRig                        <https://github.com/xmrig>, <support@xmrig.com>
  *
  *   This program is free software: you can redistribute it and/or modify
@@ -82,7 +83,7 @@ static inline std::string sysfs_prefix(const PciTopology &topology)
 
     for (uint32_t i = 1; i < 10; ++i) {
         const std::string prefix = path + std::to_string(i) + "/";
-        if (sysfs_is_amdgpu(prefix + "name") && sysfs_read(prefix + "freq1_input")) {
+        if (sysfs_is_amdgpu(prefix + "name") && (sysfs_read(prefix + "temp1_input") || sysfs_read(prefix + "power1_average"))) {
             return prefix;
         }
     }
@@ -116,6 +117,7 @@ void xmrig::AdlLib::close()
 }
 
 
+// https://dri.freedesktop.org/docs/drm/gpu/amdgpu.html#gpu-power-thermal-controls-and-monitoring
 AdlHealth xmrig::AdlLib::health(const OclDevice &device)
 {
     if (!isReady() || device.vendorId() != OCL_VENDOR_AMD) {
@@ -134,6 +136,10 @@ AdlHealth xmrig::AdlLib::health(const OclDevice &device)
     health.rpm          = sysfs_read(prefix + "fan1_input");
     health.temperature  = sysfs_read(prefix + "temp2_input") / 1000;
 
+    if (!health.temperature) {
+        health.temperature = sysfs_read(prefix + "temp1_input") / 1000;
+    }
+
     return health;
 }
 

src/base/net/stratum/NetworkState.cpp

@@ -25,6 +25,7 @@
 
 #include "base/net/stratum/NetworkState.h"
 #include "3rdparty/rapidjson/document.h"
+#include "base/io/log/Log.h"
 #include "base/kernel/interfaces/IClient.h"
 #include "base/kernel/interfaces/IStrategy.h"
 #include "base/net/stratum/Job.h"
@@ -40,6 +41,80 @@
 
 
 
+namespace xmrig {
+
+
+inline static void printCount(uint64_t accepted, uint64_t rejected)
+{
+    float percent   = 100.0;
+    int color       = 2;
+
+    if (!accepted) {
+        percent     = 0.0;
+        color       = 1;
+    }
+    else if (rejected) {
+        percent     = static_cast<float>(accepted) / (accepted + rejected) * 100.0;
+        color       = 3;
+    }
+
+    Log::print(GREEN_BOLD(" * ") WHITE_BOLD("%-17s") CSI "1;3%dm%" PRIu64 CLEAR CSI "0;3%dm (%1.1f%%)", "accepted", color, accepted, color, percent);
+
+    if (rejected) {
+        Log::print(GREEN_BOLD(" * ") WHITE_BOLD("%-17s") RED_BOLD("%" PRIu64), "rejected", rejected);
+    }
+}
+
+
+inline static void printHashes(uint64_t accepted, uint64_t hashes)
+{
+    Log::print(GREEN_BOLD(" * ") WHITE_BOLD("%-17s") CYAN_BOLD("%" PRIu64) " avg " CYAN("%1.0f"),
+               "pool-side hashes", hashes, static_cast<double>(hashes) / accepted);
+}
+
+
+inline static void printAvgTime(uint64_t time)
+{
+    Log::print(GREEN_BOLD(" * ") WHITE_BOLD("%-17s") CSI "1;3%dm%1.1fs", "avg result time", (time < 10000 ? 3 : 2), time / 1000.0);
+}
+
+
+static void printDiff(uint64_t diff)
+{
+    Log::print(GREEN_BOLD(" * ") WHITE_BOLD("%-17s") CYAN_BOLD("%s"), "difficulty", NetworkState::humanDiff(diff).c_str());
+}
+
+
+inline static void printDiff(size_t i, uint64_t diff, uint64_t hashes)
+{
+    if (!diff) {
+        return;
+    }
+
+    const double effort = static_cast<double>(hashes) / diff * 100.0;
+    const double target = (i + 1) * 100.0;
+    const int color     = effort > (target + 100.0) ? 1 : (effort > target ? 3 : 2);
+
+    Log::print("%3zu | %10s | " CSI "0;3%dm%8.2f" CLEAR " |", i + 1, NetworkState::humanDiff(diff).c_str(), color, effort);
+}
+
+
+inline static void printLatency(uint32_t latency)
+{
+    if (!latency) {
+        return;
+    }
+
+    const int color = latency < 100 ? 2 : (latency > 500 ? 1 : 3);
+
+    Log::print(GREEN_BOLD(" * ") WHITE_BOLD("%-17s") CSI "1;3%dm%ums", "ping time", color, latency);
+}
+
+
+} // namespace xmrig
+
+
+
 xmrig::NetworkState::NetworkState(IStrategyListener *listener) : StrategyProxy(listener)
 {
 }
@@ -54,7 +129,8 @@ rapidjson::Value xmrig::NetworkState::getConnection(rapidjson::Document &doc, in
     Value connection(kObjectType);
     connection.AddMember("pool",            StringRef(m_pool), allocator);
     connection.AddMember("ip",              m_ip.toJSON(), allocator);
-    connection.AddMember("uptime",          connectionTime(), allocator);
+    connection.AddMember("uptime",          connectionTime() / 1000, allocator);
+    connection.AddMember("uptime_ms",       connectionTime(), allocator);
     connection.AddMember("ping",            latency(), allocator);
     connection.AddMember("failures",        m_failures, allocator);
     connection.AddMember("tls",             m_tls.toJSON(), allocator);
@@ -64,7 +140,8 @@ rapidjson::Value xmrig::NetworkState::getConnection(rapidjson::Document &doc, in
     connection.AddMember("diff",            m_diff, allocator);
     connection.AddMember("accepted",        m_accepted, allocator);
     connection.AddMember("rejected",        m_rejected, allocator);
-    connection.AddMember("avg_time",        avgTime(), allocator);
+    connection.AddMember("avg_time",        avgTime() / 1000, allocator);
+    connection.AddMember("avg_time_ms",     avgTime(), allocator);
     connection.AddMember("hashes_total",    m_hashes, allocator);
 
     if (version == 1) {
@@ -85,11 +162,14 @@ rapidjson::Value xmrig::NetworkState::getResults(rapidjson::Document &doc, int v
     results.AddMember("diff_current",  m_diff, allocator);
     results.AddMember("shares_good",   m_accepted, allocator);
     results.AddMember("shares_total",  m_accepted + m_rejected, allocator);
-    results.AddMember("avg_time",      avgTime(), allocator);
+    results.AddMember("avg_time",      avgTime() / 1000, allocator);
+    results.AddMember("avg_time_ms",   avgTime(), allocator);
     results.AddMember("hashes_total",  m_hashes, allocator);
 
     Value best(kArrayType);
-    for (uint64_t i : topDiff) {
+    best.Reserve(m_topDiff.size(), allocator);
+
+    for (uint64_t i : m_topDiff) {
         best.PushBack(i, allocator);
     }
 
@@ -104,6 +184,84 @@ rapidjson::Value xmrig::NetworkState::getResults(rapidjson::Document &doc, int v
 #endif
 
 
+void xmrig::NetworkState::printConnection() const
+{
+    if (!m_active) {
+        LOG_NOTICE(YELLOW_BOLD_S "no active connection");
+
+        return;
+    }
+
+    Log::print(MAGENTA_BOLD_S " - CONNECTION");
+    Log::print(GREEN_BOLD(" * ") WHITE_BOLD("%-17s") CYAN_BOLD("%s ") BLACK_BOLD("(%s) ") GREEN_BOLD("%s"),
+               "pool address", m_pool, m_ip.data(), m_tls.isNull() ? "" : m_tls.data());
+
+    Log::print(GREEN_BOLD(" * ") WHITE_BOLD("%-17s") WHITE_BOLD("%s"), "algorithm", m_algorithm.shortName());
+    printDiff(m_diff);
+    printLatency(latency());
+    Log::print(GREEN_BOLD(" * ") WHITE_BOLD("%-17s") CYAN_BOLD("%" PRIu64 "s"), "connection time", connectionTime() / 1000);
+}
+
+
+void xmrig::NetworkState::printResults() const
+{
+    if (!m_hashes) {
+        LOG_NOTICE(YELLOW_BOLD_S "no any results yet");
+
+        return;
+    }
+
+    Log::print(MAGENTA_BOLD_S " - RESULTS");
+
+    printCount(m_accepted, m_rejected);
+    printHashes(m_accepted, m_hashes);
+    printDiff(m_diff);
+
+    if (m_active && !m_latency.empty()) {
+        printAvgTime(avgTime());
+    }
+
+    Log::print(MAGENTA_BOLD_S " - TOP 10");
+    Log::print(WHITE_BOLD_S "  # | DIFFICULTY | EFFORT %% |");
+
+    for (size_t i = 0; i < m_topDiff.size(); ++i) {
+        printDiff(i, m_topDiff[i], m_hashes);
+    }
+}
+
+
+const char *xmrig::NetworkState::scaleDiff(uint64_t &diff)
+{
+    if (diff >= 100000000000) {
+        diff /= 1000000000;
+
+        return "G";
+    }
+
+    if (diff >= 100000000) {
+        diff /= 1000000;
+
+        return "M";
+    }
+
+    if (diff >= 1000000) {
+        diff /= 1000;
+
+        return "K";
+    }
+
+    return "";
+}
+
+
+std::string xmrig::NetworkState::humanDiff(uint64_t diff)
+{
+    const char *scale = scaleDiff(diff);
+
+    return std::to_string(diff) + scale;
+}
+
+
 void xmrig::NetworkState::onActive(IStrategy *strategy, IClient *client)
 {
     snprintf(m_pool, sizeof(m_pool) - 1, "%s:%d", client->pool().host().data(), client->pool().port());
@@ -145,16 +303,6 @@ void xmrig::NetworkState::onResultAccepted(IStrategy *strategy, IClient *client,
 }
 
 
-uint32_t xmrig::NetworkState::avgTime() const
-{
-    if (m_latency.empty()) {
-        return 0;
-    }
-
-    return connectionTime() / (uint32_t)m_latency.size();
-}
-
-
 uint32_t xmrig::NetworkState::latency() const
 {
     const size_t calls = m_latency.size();
@@ -169,9 +317,19 @@ uint32_t xmrig::NetworkState::latency() const
 }
 
 
+uint64_t xmrig::NetworkState::avgTime() const
+{
+    if (m_latency.empty()) {
+        return 0;
+    }
+
+    return connectionTime() / m_latency.size();
+}
+
+
 uint64_t xmrig::NetworkState::connectionTime() const
 {
-    return m_active ? ((Chrono::steadyMSecs() - m_connectionTime) / 1000) : 0;
+    return m_active ? ((Chrono::steadyMSecs() - m_connectionTime)) : 0;
 }
 
 
@@ -185,10 +343,10 @@ void xmrig::NetworkState::add(const SubmitResult &result, const char *error)
     m_accepted++;
     m_hashes += result.diff;
 
-    const size_t ln = topDiff.size() - 1;
-    if (result.actualDiff > topDiff[ln]) {
-        topDiff[ln] = result.actualDiff;
-        std::sort(topDiff.rbegin(), topDiff.rend());
+    const size_t ln = m_topDiff.size() - 1;
+    if (result.actualDiff > m_topDiff[ln]) {
+        m_topDiff[ln] = result.actualDiff;
+        std::sort(m_topDiff.rbegin(), m_topDiff.rend());
     }
 
     m_latency.push_back(result.elapsed > 0xFFFF ? 0xFFFF : static_cast<uint16_t>(result.elapsed));

src/base/net/stratum/NetworkState.h

@@ -32,6 +32,7 @@
 
 
 #include <array>
+#include <string>
 #include <vector>
 
 
@@ -52,6 +53,12 @@ public:
     rapidjson::Value getResults(rapidjson::Document &doc, int version) const;
 #   endif
 
+    void printConnection() const;
+    void printResults() const;
+
+    static const char *scaleDiff(uint64_t &diff);
+    static std::string humanDiff(uint64_t diff);
+
 protected:
     void onActive(IStrategy *strategy, IClient *client) override;
     void onJob(IStrategy *strategy, IClient *client, const Job &job, const rapidjson::Value &params) override;
@@ -59,8 +66,8 @@ protected:
     void onResultAccepted(IStrategy *strategy, IClient *client, const SubmitResult &result, const char *error) override;
 
 private:
-    uint32_t avgTime() const;
     uint32_t latency() const;
+    uint64_t avgTime() const;
     uint64_t connectionTime() const;
     void add(const SubmitResult &result, const char *error);
     void stop();
@@ -68,7 +75,7 @@ private:
     Algorithm m_algorithm;
     bool m_active               = false;
     char m_pool[256]{};
-    std::array<uint64_t, 10> topDiff { { } };
+    std::array<uint64_t, 10> m_topDiff { { } };
     std::vector<uint16_t> m_latency;
     String m_fingerprint;
     String m_ip;

src/core/Controller.cpp

@@ -98,3 +98,10 @@ xmrig::Network *xmrig::Controller::network() const
 
     return m_network;
 }
+
+
+void xmrig::Controller::execCommand(char command)
+{
+    miner()->execCommand(command);
+    network()->execCommand(command);
+}

src/core/Controller.h

@@ -52,6 +52,7 @@ public:
 
     Miner *miner() const;
     Network *network() const;
+    void execCommand(char command);
 
 private:
     Miner *m_miner     = nullptr;

src/crypto/kawpow/KPCache.cpp

@@ -92,9 +92,9 @@ bool KPCache::init(uint32_t epoch)
             const uint32_t b = (cache_nodes * (i + 1)) / n;
 
             threads.emplace_back([this, a, b, cache_nodes, &cache]() {
-                for (uint32_t j = a; j < b; ++j) {
-                    ethash_calculate_dag_item_opt(((node*)m_DAGCache.data()) + j, j, num_dataset_parents, &cache);
-                }
+                uint32_t j = a;
+                for (; j + 4 <= b; j += 4) ethash_calculate_dag_item4_opt(((node*)m_DAGCache.data()) + j, j, num_dataset_parents, &cache);
+                for (; j < b; ++j) ethash_calculate_dag_item_opt(((node*)m_DAGCache.data()) + j, j, num_dataset_parents, &cache);
             });
         }
 

src/crypto/kawpow/KPHash.cpp

@@ -285,10 +285,7 @@ void KPHash::calculate(const KPCache& light_cache, uint32_t block_height, const
         uint32_t item_index = (mix[r % LANES][0] % num_items) * 4;
 
         node item[4];
-        ethash_calculate_dag_item_opt(item + 0, item_index + 0, KPCache::num_dataset_parents, &cache);
-        ethash_calculate_dag_item_opt(item + 1, item_index + 1, KPCache::num_dataset_parents, &cache);
-        ethash_calculate_dag_item_opt(item + 2, item_index + 2, KPCache::num_dataset_parents, &cache);
-        ethash_calculate_dag_item_opt(item + 3, item_index + 3, KPCache::num_dataset_parents, &cache);
+        ethash_calculate_dag_item4_opt(item, item_index, KPCache::num_dataset_parents, &cache);
 
         uint32_t dst_counter = 0;
         uint32_t src_counter = 0;

src/crypto/randomx/asm/program_read_dataset_ryzen.inc

@@ -1,13 +1,12 @@
 	mov rcx, rbp                       ;# ecx = ma
 	shr rcx, 32
 	and ecx, RANDOMX_DATASET_BASE_MASK
+	xor r8, qword ptr [rdi+rcx]
 	xor rbp, rax                       ;# modify "mx"
-	mov rax, qword ptr [rdi+rcx]
 	mov edx, ebp                       ;# edx = mx
 	and edx, RANDOMX_DATASET_BASE_MASK
 	prefetchnta byte ptr [rdi+rdx]
 	ror rbp, 32                        ;# swap "ma" and "mx"
-	xor r8,  rax
 	xor r9,  qword ptr [rdi+rcx+8]
 	xor r10, qword ptr [rdi+rcx+16]
 	xor r11, qword ptr [rdi+rcx+24]

src/crypto/randomx/randomx.h

@@ -121,7 +121,7 @@ struct RandomX_ConfigurationBase
 	uint8_t codeShhPrefetchTweaked[20];
 	uint8_t codeReadDatasetTweaked[64];
 	uint32_t codeReadDatasetTweakedSize;
-	uint8_t codeReadDatasetRyzenTweaked[76];
+	uint8_t codeReadDatasetRyzenTweaked[72];
 	uint32_t codeReadDatasetRyzenTweakedSize;
 	uint8_t codeReadDatasetLightSshInitTweaked[68];
 	uint8_t codePrefetchScratchpadTweaked[32];

src/crypto/rx/Rx_linux.cpp

@@ -228,6 +228,9 @@ void xmrig::Rx::msrInit(const RxConfig &config)
     if (wrmsr(preset, config.rdmsr())) {
         LOG_NOTICE(CLEAR "%s" GREEN_BOLD_S "register values for \"%s\" preset has been set successfully" BLACK_BOLD(" (%" PRIu64 " ms)"), tag, config.msrPresetName(), Chrono::steadyMSecs() - ts);
     }
+    else {
+        LOG_ERR(CLEAR "%s" RED_BOLD_S "FAILED TO APPLY MSR MOD, HASHRATE WILL BE LOW", tag);
+    }
 }
 
 

src/crypto/rx/Rx_win.cpp

@@ -361,6 +361,9 @@ void xmrig::Rx::msrInit(const RxConfig &config)
     if (wrmsr(preset, config.rdmsr())) {
         LOG_NOTICE(CLEAR "%s" GREEN_BOLD_S "register values for \"%s\" preset has been set successfully" BLACK_BOLD(" (%" PRIu64 " ms)"), tag, config.msrPresetName(), Chrono::steadyMSecs() - ts);
     }
+    else {
+        LOG_ERR(CLEAR "%s" RED_BOLD_S "FAILED TO APPLY MSR MOD, HASHRATE WILL BE LOW", tag);
+    }
 }
 
 

src/net/Network.cpp

@@ -98,6 +98,25 @@ void xmrig::Network::connect()
 }
 
 
+void xmrig::Network::execCommand(char command)
+{
+    switch (command) {
+    case 's':
+    case 'S':
+        m_state->printResults();
+        break;
+
+    case 'c':
+    case 'C':
+        m_state->printConnection();
+        break;
+
+    default:
+        break;
+    }
+}
+
+
 void xmrig::Network::onActive(IStrategy *strategy, IClient *client)
 {
     if (m_donate && m_donate == strategy) {
@@ -192,24 +211,10 @@ void xmrig::Network::onPause(IStrategy *strategy)
 }
 
 
-static void scale_diff(uint64_t& diff, const char* &scale)
-{
-    if (diff >= 100000000) {
-        diff /= 1000000;
-        scale = "M";
-    }
-    else if (diff >= 1000000) {
-        diff /= 1000;
-        scale = "K";
-    }
-}
-
-
 void xmrig::Network::onResultAccepted(IStrategy *, IClient *, const SubmitResult &result, const char *error)
 {
-    uint64_t diff = result.diff;
-    const char* scale = "";
-    scale_diff(diff, scale);
+    uint64_t diff     = result.diff;
+    const char *scale = NetworkState::scaleDiff(diff);
 
     if (error) {
         LOG_INFO("%s " RED_BOLD("rejected") " (%" PRId64 "/%" PRId64 ") diff " WHITE_BOLD("%" PRIu64 "%s") " " RED("\"%s\"") " " BLACK_BOLD("(%" PRIu64 " ms)"),
@@ -247,9 +252,8 @@ void xmrig::Network::onRequest(IApiRequest &request)
 
 void xmrig::Network::setJob(IClient *client, const Job &job, bool donate)
 {
-    uint64_t diff = job.diff();
-    const char* scale = "";
-    scale_diff(diff, scale);
+    uint64_t diff     = job.diff();;
+    const char *scale = NetworkState::scaleDiff(diff);
 
     if (job.height()) {
         LOG_INFO("%s " MAGENTA_BOLD("new job") " from " WHITE_BOLD("%s:%d") " diff " WHITE_BOLD("%" PRIu64 "%s") " algo " WHITE_BOLD("%s") " height " WHITE_BOLD("%" PRIu64),

src/net/Network.h

@@ -58,6 +58,7 @@ public:
     inline IStrategy *strategy() const { return m_strategy; }
 
     void connect();
+    void execCommand(char command);
 
 protected:
     inline void onTimer(const Timer *) override { tick(); }

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.2"
+#define APP_VERSION   "6.2.3"
 #define APP_DOMAIN    "xmrig.com"
 #define APP_SITE      "www.xmrig.com"
 #define APP_COPYRIGHT "Copyright (C) 2016-2020 xmrig.com"
@@ -36,7 +36,7 @@
 
 #define APP_VER_MAJOR  6
 #define APP_VER_MINOR  2
-#define APP_VER_PATCH  2
+#define APP_VER_PATCH  3
 
 #ifdef _MSC_VER
 #   if (_MSC_VER >= 1920)