Merge a776ebf394 into 98c775703e

CHANGELOG.md

@@ -1,11 +1,3 @@
-# v6.22.0
-- [#2411](https://github.com/xmrig/xmrig/pull/2411) Added support for [Yada](https://yadacoin.io/) (`rx/yada` algorithm).
-- [#3492](https://github.com/xmrig/xmrig/pull/3492) Fixed `--background` option on Unix systems.
-- [#3518](https://github.com/xmrig/xmrig/pull/3518) Possible fix for corrupted API output in rare cases.
-- [#3522](https://github.com/xmrig/xmrig/pull/3522) Removed `rx/keva` algorithm.
-- [#3525](https://github.com/xmrig/xmrig/pull/3525) Added Zen5 detection.
-- [#3528](https://github.com/xmrig/xmrig/pull/3528) Added `rx/yada` OpenCL support.
-
 # v6.21.3
 - [#3462](https://github.com/xmrig/xmrig/pull/3462) RandomX: correct memcpy size for JIT initialization.
 

CMakeLists.txt

@@ -14,7 +14,9 @@ option(WITH_HTTP            "Enable HTTP protocol support (client/server)" ON)
 option(WITH_DEBUG_LOG       "Enable debug log output" OFF)
 option(WITH_TLS             "Enable OpenSSL support" ON)
 option(WITH_ASM             "Enable ASM PoW implementations" ON)
-option(WITH_MSR             "Enable MSR mod & 1st-gen Ryzen fix" ON)
+option(WITH_ASM_AMD         "Enable ASM for AMD processors" ON)
+option(WITH_MSR             "Enable MSR mod" ON)
+option(WITH_MSR_ZEN         "Enable MSR mod for AMD Zen-based processors" ON)
 option(WITH_ENV_VARS        "Enable environment variables support in config file" ON)
 option(WITH_EMBEDDED_CONFIG "Enable internal embedded JSON config" OFF)
 option(WITH_OPENCL          "Enable OpenCL backend" ON)

cmake/asm.cmake

@@ -44,9 +44,17 @@ if (WITH_ASM AND NOT XMRIG_ARM AND CMAKE_SIZEOF_VOID_P EQUAL 8)
     set_property(TARGET ${XMRIG_ASM_LIBRARY} PROPERTY LINKER_LANGUAGE C)
 
     add_definitions(/DXMRIG_FEATURE_ASM)
+    if (WITH_ASM_AMD)
+        add_definitions(/DXMRIG_FEATURE_ASM_AMD)
+        message("-- WITH_ASM=ON (+amd)")
+    else()
+        message("-- WITH_ASM=ON (-amd)")
+    endif()
 else()
     set(XMRIG_ASM_SOURCES "")
     set(XMRIG_ASM_LIBRARY "")
 
     remove_definitions(/DXMRIG_FEATURE_ASM)
+    remove_definitions(/DXMRIG_FEATURE_ASM_AMD)
+    message("-- WITH_ASM=OFF")
 endif()

cmake/randomx.cmake

@@ -104,8 +104,13 @@ if (WITH_RANDOMX)
 
     if (WITH_MSR AND NOT XMRIG_ARM AND CMAKE_SIZEOF_VOID_P EQUAL 8 AND (XMRIG_OS_WIN OR XMRIG_OS_LINUX))
         add_definitions(/DXMRIG_FEATURE_MSR)
-        add_definitions(/DXMRIG_FIX_RYZEN)
+        if (WITH_MSR_ZEN)
-        message("-- WITH_MSR=ON")
+            add_definitions(/DXMRIG_FIX_RYZEN)
+            message("-- WITH_MSR=ON (+zen)")
+        else()
+            remove_definitions(/DXMRIG_FIX_RYZEN)
+            message("-- WITH_MSR=ON (-zen)")
+        endif()
 
         if (XMRIG_OS_WIN)
             list(APPEND SOURCES_CRYPTO

src/backend/opencl/OclWorker.cpp

@@ -170,7 +170,7 @@ void xmrig::OclWorker::start()
             const uint64_t t = Chrono::steadyMSecs();
 
             try {
-                m_runner->run(readUnaligned(m_job.nonce()), m_job.nonceOffset(), results);
+                m_runner->run(readUnaligned(m_job.nonce()), results);
             }
             catch (std::exception &ex) {
                 printError(id(), ex.what());

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

@@ -23,7 +23,6 @@
 #define ALGO_RX_ARQMA       0x72121061
 #define ALGO_RX_SFX         0x72151273
 #define ALGO_RX_GRAFT       0x72151267
-#define ALGO_RX_YADA        0x72151279
 #define ALGO_AR2_CHUKWA     0x61130000
 #define ALGO_AR2_CHUKWA_V2  0x61140000
 #define ALGO_AR2_WRKZ       0x61120000

src/backend/opencl/cl/cn/cryptonight_cl.h

@@ -2,7 +2,7 @@
 
 namespace xmrig {
 
-static const char cryptonight_cl[61447] = {
+static const char cryptonight_cl[61415] = {
     0x23,0x69,0x66,0x64,0x65,0x66,0x20,0x53,0x54,0x41,0x54,0x49,0x43,0x0a,0x23,0x75,0x6e,0x64,0x65,0x66,0x20,0x53,0x54,0x41,0x54,0x49,0x43,0x0a,0x23,0x65,0x6e,0x64,
     0x69,0x66,0x0a,0x23,0x69,0x66,0x64,0x65,0x66,0x20,0x63,0x6c,0x5f,0x61,0x6d,0x64,0x5f,0x6d,0x65,0x64,0x69,0x61,0x5f,0x6f,0x70,0x73,0x0a,0x23,0x64,0x65,0x66,0x69,
     0x6e,0x65,0x20,0x53,0x54,0x41,0x54,0x49,0x43,0x20,0x73,0x74,0x61,0x74,0x69,0x63,0x0a,0x23,0x65,0x6c,0x73,0x65,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x53,
@@ -35,8 +35,7 @@ static const char cryptonight_cl[61447] = {
     0x34,0x31,0x31,0x37,0x37,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x41,0x52,0x51,0x4d,0x41,0x20,0x30,0x78,0x37,0x32,
     0x31,0x32,0x31,0x30,0x36,0x31,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x53,0x46,0x58,0x20,0x30,0x78,0x37,0x32,0x31,
     0x35,0x31,0x32,0x37,0x33,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x47,0x52,0x41,0x46,0x54,0x20,0x30,0x78,0x37,0x32,
-    0x31,0x35,0x31,0x32,0x36,0x37,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x59,0x41,0x44,0x41,0x20,0x30,0x78,0x37,0x32,
+    0x31,0x35,0x31,0x32,0x36,0x37,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,0x30,
-    0x31,0x35,0x31,0x32,0x37,0x39,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,0x30,
     0x78,0x36,0x31,0x31,0x33,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,
     0x41,0x5f,0x56,0x32,0x20,0x30,0x78,0x36,0x31,0x31,0x34,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,
     0x5f,0x57,0x52,0x4b,0x5a,0x20,0x30,0x78,0x36,0x31,0x31,0x32,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x4b,0x41,

src/backend/opencl/cl/rx/blake2b.cl

@@ -225,110 +225,6 @@ __kernel void blake2b_initial_hash_double(__global void *out, __global const voi
 	t[7] = hash[7];
 }
 
-void blake2b_512_process_big_block(ulong *out, __global const ulong* in, uint in_len, uint out_len, uint nonce, uint nonce_offset)
-{
-	ulong h[8] = { iv0 ^ (0x01010000u | out_len), iv1, iv2, iv3, iv4, iv5, iv6, iv7 };
-
-	for (uint t = 128; t < in_len; t += 128, in += 16) {
-		ulong m[16] = { in[0], in[1], in[2], in[3], in[4], in[5], in[6], in[7], in[8], in[9], in[10], in[11],  in[12], in[13], in[14], in[15] };
-
-		const uint k0 = (nonce_offset + 0) - (t - 128);
-		const uint k1 = (nonce_offset + 1) - (t - 128);
-		const uint k2 = (nonce_offset + 2) - (t - 128);
-		const uint k3 = (nonce_offset + 3) - (t - 128);
-
-		if (k0 < 128) m[k0 / 8] |= (ulong)((nonce >>  0) & 255) << ((k0 % 8) * 8);
-		if (k1 < 128) m[k1 / 8] |= (ulong)((nonce >>  8) & 255) << ((k1 % 8) * 8);
-		if (k2 < 128) m[k2 / 8] |= (ulong)((nonce >> 16) & 255) << ((k2 % 8) * 8);
-		if (k3 < 128) m[k3 / 8] |= (ulong)((nonce >> 24) & 255) << ((k3 % 8) * 8);
-
-		ulong v[16] = {  h[0],  h[1],  h[2],  h[3],  h[4],  h[5],  h[6],  h[7],   iv0,   iv1,    iv2,    iv3, iv4 ^ t,    iv5,    iv6,    iv7 };
-
-		BLAKE2B_ROUNDS();
-
-		h[0] ^= v[0] ^ v[ 8];
-		h[1] ^= v[1] ^ v[ 9];
-		h[2] ^= v[2] ^ v[10];
-		h[3] ^= v[3] ^ v[11];
-		h[4] ^= v[4] ^ v[12];
-		h[5] ^= v[5] ^ v[13];
-		h[6] ^= v[6] ^ v[14];
-		h[7] ^= v[7] ^ v[15];
-	}
-
-	uint k = in_len & 127;
-	if (k == 0) k = 128;
-
-	ulong m[16] = {
-		(k >   0) ? in[ 0] : 0,
-		(k >   8) ? in[ 1] : 0,
-		(k >  16) ? in[ 2] : 0,
-		(k >  24) ? in[ 3] : 0,
-		(k >  32) ? in[ 4] : 0,
-		(k >  40) ? in[ 5] : 0,
-		(k >  48) ? in[ 6] : 0,
-		(k >  56) ? in[ 7] : 0,
-		(k >  64) ? in[ 8] : 0,
-		(k >  72) ? in[ 9] : 0,
-		(k >  80) ? in[10] : 0,
-		(k >  88) ? in[11] : 0,
-		(k >  96) ? in[12] : 0,
-		(k > 104) ? in[13] : 0,
-		(k > 112) ? in[14] : 0,
-		(k > 120) ? in[15] : 0
-	};
-
-	const uint t = in_len - k;
-
-	const uint k0 = nonce_offset + 0 - t;
-	const uint k1 = nonce_offset + 1 - t;
-	const uint k2 = nonce_offset + 2 - t;
-	const uint k3 = nonce_offset + 3 - t;
-
-	if (k0 < k) m[k0 / 8] |= (ulong)((nonce >>  0) & 255) << ((k0 % 8) * 8);
-	if (k1 < k) m[k1 / 8] |= (ulong)((nonce >>  8) & 255) << ((k1 % 8) * 8);
-	if (k2 < k) m[k2 / 8] |= (ulong)((nonce >> 16) & 255) << ((k2 % 8) * 8);
-	if (k3 < k) m[k3 / 8] |= (ulong)((nonce >> 24) & 255) << ((k3 % 8) * 8);
-
-	if (k % 8) {
-		m[k / 8] &= (ulong)(-1) >> (64 - (k % 8) * 8);
-	}
-
-	ulong v[16] = { h[0], h[1], h[2], h[3], h[4], h[5], h[6], h[7], iv0, iv1, iv2, iv3, iv4 ^ in_len, iv5, ~iv6, iv7 };
-
-	BLAKE2B_ROUNDS();
-
-	if (out_len >  0) out[0] = h[0] ^ v[0] ^ v[8];
-	if (out_len >  8) out[1] = h[1] ^ v[1] ^ v[9];
-	if (out_len > 16) out[2] = h[2] ^ v[2] ^ v[10];
-	if (out_len > 24) out[3] = h[3] ^ v[3] ^ v[11];
-	if (out_len > 32) out[4] = h[4] ^ v[4] ^ v[12];
-	if (out_len > 40) out[5] = h[5] ^ v[5] ^ v[13];
-	if (out_len > 48) out[6] = h[6] ^ v[6] ^ v[14];
-	if (out_len > 56) out[7] = h[7] ^ v[7] ^ v[15];
-}
-
-__attribute__((reqd_work_group_size(64, 1, 1)))
-__kernel void blake2b_initial_hash_big(__global void *out, __global const void* blockTemplate, uint blockTemplateSize, uint start_nonce, uint nonce_offset)
-{
-	const uint global_index = get_global_id(0);
-
-	__global const ulong* p = (__global const ulong*) blockTemplate;
-
-	ulong hash[8];
-	blake2b_512_process_big_block(hash, p, blockTemplateSize, 64, start_nonce + global_index, nonce_offset);
-
-	__global ulong* t = ((__global ulong*) out) + global_index * 8;
-	t[0] = hash[0];
-	t[1] = hash[1];
-	t[2] = hash[2];
-	t[3] = hash[3];
-	t[4] = hash[4];
-	t[5] = hash[5];
-	t[6] = hash[6];
-	t[7] = hash[7];
-}
-
 #define in_len 256
 
 #define out_len 32

src/backend/opencl/cl/rx/randomx.cl

@@ -1,6 +1,6 @@
 #include "../cn/algorithm.cl"
 
-#if ((ALGO == ALGO_RX_0) || (ALGO == ALGO_RX_YADA))
+#if (ALGO == ALGO_RX_0)
 #include "randomx_constants_monero.h"
 #elif (ALGO == ALGO_RX_WOW)
 #include "randomx_constants_wow.h"

src/backend/opencl/interfaces/IOclRunner.h

@@ -64,7 +64,7 @@ public:
     virtual uint32_t deviceIndex() const                    = 0;
     virtual void build()                                    = 0;
     virtual void init()                                     = 0;
-    virtual void run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput)  = 0;
+    virtual void run(uint32_t nonce, uint32_t *hashOutput)  = 0;
     virtual void set(const Job &job, uint8_t *blob)         = 0;
     virtual void jobEarlyNotification(const Job&)           = 0;
 

src/backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.cpp

@@ -1,59 +0,0 @@
-/* XMRig
- * Copyright 2010      Jeff Garzik <jgarzik@pobox.com>
- * Copyright 2012-2014 pooler      <pooler@litecoinpool.org>
- * Copyright 2014      Lucas Jones <https://github.com/lucasjones>
- * Copyright 2014-2016 Wolf9466    <https://github.com/OhGodAPet>
- * Copyright 2016      Jay D Dee   <jayddee246@gmail.com>
- * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
- * Copyright 2018-2019 SChernykh   <https://github.com/SChernykh>
- * Copyright 2016-2019 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
- *
- *   This program is free software: you can redistribute it and/or modify
- *   it under the terms of the GNU General Public License as published by
- *   the Free Software Foundation, either version 3 of the License, or
- *   (at your option) any later version.
- *
- *   This program is distributed in the hope that it will be useful,
- *   but WITHOUT ANY WARRANTY; without even the implied warranty of
- *   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
- *   GNU General Public License for more details.
- *
- *   You should have received a copy of the GNU General Public License
- *   along with this program. If not, see <http://www.gnu.org/licenses/>.
- */
-
-
-#include "backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.h"
-#include "backend/opencl/wrappers/OclLib.h"
-
-
-void xmrig::Blake2bInitialHashBigKernel::enqueue(cl_command_queue queue, size_t threads)
-{
-    const size_t gthreads        = threads;
-    static const size_t lthreads = 64;
-
-    enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
-}
-
-
-// __kernel void blake2b_initial_hash_double(__global void *out, __global const void* blockTemplate, uint blockTemplateSize, uint start_nonce)
-void xmrig::Blake2bInitialHashBigKernel::setArgs(cl_mem out, cl_mem blockTemplate)
-{
-    setArg(0, sizeof(cl_mem), &out);
-    setArg(1, sizeof(cl_mem), &blockTemplate);
-}
-
-
-void xmrig::Blake2bInitialHashBigKernel::setBlobSize(size_t size)
-{
-    const uint32_t s = size;
-
-    setArg(2, sizeof(uint32_t), &s);
-}
-
-
-void xmrig::Blake2bInitialHashBigKernel::setNonce(uint32_t nonce, uint32_t nonce_offset)
-{
-    setArg(3, sizeof(uint32_t), &nonce);
-    setArg(4, sizeof(uint32_t), &nonce_offset);
-}

src/backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.h

@@ -1,50 +0,0 @@
-/* XMRig
- * Copyright 2010      Jeff Garzik <jgarzik@pobox.com>
- * Copyright 2012-2014 pooler      <pooler@litecoinpool.org>
- * Copyright 2014      Lucas Jones <https://github.com/lucasjones>
- * Copyright 2014-2016 Wolf9466    <https://github.com/OhGodAPet>
- * Copyright 2016      Jay D Dee   <jayddee246@gmail.com>
- * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
- * Copyright 2018-2019 SChernykh   <https://github.com/SChernykh>
- * Copyright 2016-2019 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
- *
- *   This program is free software: you can redistribute it and/or modify
- *   it under the terms of the GNU General Public License as published by
- *   the Free Software Foundation, either version 3 of the License, or
- *   (at your option) any later version.
- *
- *   This program is distributed in the hope that it will be useful,
- *   but WITHOUT ANY WARRANTY; without even the implied warranty of
- *   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
- *   GNU General Public License for more details.
- *
- *   You should have received a copy of the GNU General Public License
- *   along with this program. If not, see <http://www.gnu.org/licenses/>.
- */
-
-#ifndef XMRIG_BLAKE2BINITIALHASHBIGKERNEL_H
-#define XMRIG_BLAKE2BINITIALHASHBIGKERNEL_H
-
-
-#include "backend/opencl/wrappers/OclKernel.h"
-
-
-namespace xmrig {
-
-
-class Blake2bInitialHashBigKernel : public OclKernel
-{
-public:
-    inline Blake2bInitialHashBigKernel(cl_program program) : OclKernel(program, "blake2b_initial_hash_big") {}
-
-    void enqueue(cl_command_queue queue, size_t threads);
-    void setArgs(cl_mem out, cl_mem blockTemplate);
-    void setBlobSize(size_t size);
-    void setNonce(uint32_t nonce, uint32_t nonce_offset);
-};
-
-
-} // namespace xmrig
-
-
-#endif /* XMRIG_BLAKE2BINITIALHASHBIGKERNEL_H */

src/backend/opencl/opencl.cmake

@@ -80,7 +80,6 @@ if (WITH_OPENCL)
     if (WITH_RANDOMX)
         list(APPEND HEADERS_BACKEND_OPENCL
              src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h
-             src/backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.h
              src/backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.h
              src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.h
              src/backend/opencl/kernels/rx/ExecuteVmKernel.h
@@ -98,7 +97,6 @@ if (WITH_OPENCL)
         list(APPEND SOURCES_BACKEND_OPENCL
              src/backend/opencl/generators/ocl_generic_rx_generator.cpp
              src/backend/opencl/kernels/rx/Blake2bHashRegistersKernel.cpp
-             src/backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.cpp
              src/backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.cpp
              src/backend/opencl/kernels/rx/Blake2bInitialHashKernel.cpp
              src/backend/opencl/kernels/rx/ExecuteVmKernel.cpp

src/backend/opencl/runners/OclCnRunner.cpp

@@ -87,7 +87,7 @@ size_t xmrig::OclCnRunner::bufferSize() const
 }
 
 
-void xmrig::OclCnRunner::run(uint32_t nonce, uint32_t /*nonce_offset*/, uint32_t *hashOutput)
+void xmrig::OclCnRunner::run(uint32_t nonce, uint32_t *hashOutput)
 {
     static const cl_uint zero = 0;
 

src/backend/opencl/runners/OclCnRunner.h

@@ -42,7 +42,7 @@ public:
 
 protected:
     size_t bufferSize() const override;
-    void run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput) override;
+    void run(uint32_t nonce, uint32_t *hashOutput) override;
     void set(const Job &job, uint8_t *blob) override;
     void build() override;
     void init() override;

src/backend/opencl/runners/OclKawPowRunner.cpp

@@ -75,7 +75,7 @@ OclKawPowRunner::~OclKawPowRunner()
 }
 
 
-void OclKawPowRunner::run(uint32_t nonce, uint32_t /*nonce_offset*/, uint32_t *hashOutput)
+void OclKawPowRunner::run(uint32_t nonce, uint32_t *hashOutput)
 {
     const size_t local_work_size = m_workGroupSize;
     const size_t global_work_offset = nonce;

src/backend/opencl/runners/OclKawPowRunner.h

@@ -40,7 +40,7 @@ public:
     ~OclKawPowRunner() override;
 
 protected:
-    void run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput) override;
+    void run(uint32_t nonce, uint32_t *hashOutput) override;
     void set(const Job &job, uint8_t *blob) override;
     void build() override;
     void init() override;

src/backend/opencl/runners/OclRxBaseRunner.cpp

@@ -26,7 +26,6 @@
 #include "backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h"
 #include "backend/opencl/kernels/rx/Blake2bInitialHashKernel.h"
 #include "backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.h"
-#include "backend/opencl/kernels/rx/Blake2bInitialHashBigKernel.h"
 #include "backend/opencl/kernels/rx/FillAesKernel.h"
 #include "backend/opencl/kernels/rx/FindSharesKernel.h"
 #include "backend/opencl/kernels/rx/HashAesKernel.h"
@@ -74,7 +73,6 @@ xmrig::OclRxBaseRunner::~OclRxBaseRunner()
     delete m_hashAes1Rx4;
     delete m_blake2b_initial_hash;
     delete m_blake2b_initial_hash_double;
-    delete m_blake2b_initial_hash_big;
     delete m_blake2b_hash_registers_32;
     delete m_blake2b_hash_registers_64;
     delete m_find_shares;
@@ -87,7 +85,7 @@ xmrig::OclRxBaseRunner::~OclRxBaseRunner()
 }
 
 
-void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput)
+void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
 {
     static const uint32_t zero = 0;
 
@@ -98,7 +96,8 @@ void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t nonce_offset, uint32_t
         m_blake2b_initial_hash_double->setNonce(nonce);
     }
     else {
-        m_blake2b_initial_hash_big->setNonce(nonce, nonce_offset);
+        hashOutput[0xFF] = 0;
+        return;
     }
 
     m_find_shares->setNonce(nonce);
@@ -108,11 +107,8 @@ void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t nonce_offset, uint32_t
     if (m_jobSize <= 128) {
         m_blake2b_initial_hash->enqueue(m_queue, m_intensity);
     }
-    else if (m_jobSize <= 256) {
-        m_blake2b_initial_hash_double->enqueue(m_queue, m_intensity);
-    }
     else {
-        m_blake2b_initial_hash_big->enqueue(m_queue, m_intensity);
+        m_blake2b_initial_hash_double->enqueue(m_queue, m_intensity);
     }
 
     m_fillAes1Rx4_scratchpad->enqueue(m_queue, m_intensity);
@@ -154,15 +150,12 @@ void xmrig::OclRxBaseRunner::set(const Job &job, uint8_t *blob)
         memset(blob + job.size(), 0, Job::kMaxBlobSize - job.size());
     }
 
-    memset(blob + job.nonceOffset(), 0, job.nonceSize());
-
     enqueueWriteBuffer(m_input, CL_TRUE, 0, Job::kMaxBlobSize, blob);
 
     m_jobSize = job.size();
 
     m_blake2b_initial_hash->setBlobSize(job.size());
     m_blake2b_initial_hash_double->setBlobSize(job.size());
-    m_blake2b_initial_hash_big->setBlobSize(job.size());
 
     m_find_shares->setTarget(job.target());
 }
@@ -198,9 +191,6 @@ void xmrig::OclRxBaseRunner::build()
     m_blake2b_initial_hash_double = new Blake2bInitialHashDoubleKernel(m_program);
     m_blake2b_initial_hash_double->setArgs(m_hashes, m_input);
 
-    m_blake2b_initial_hash_big = new Blake2bInitialHashBigKernel(m_program);
-    m_blake2b_initial_hash_big->setArgs(m_hashes, m_input);
-
     m_blake2b_hash_registers_32 = new Blake2bHashRegistersKernel(m_program, "blake2b_hash_registers_32");
     m_blake2b_hash_registers_64 = new Blake2bHashRegistersKernel(m_program, "blake2b_hash_registers_64");
 

src/backend/opencl/runners/OclRxBaseRunner.h

@@ -36,7 +36,6 @@ namespace xmrig {
 class Blake2bHashRegistersKernel;
 class Blake2bInitialHashKernel;
 class Blake2bInitialHashDoubleKernel;
-class Blake2bInitialHashBigKernel;
 class FillAesKernel;
 class FindSharesKernel;
 class HashAesKernel;
@@ -54,7 +53,7 @@ protected:
     size_t bufferSize() const override;
     void build() override;
     void init() override;
-    void run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput) override;
+    void run(uint32_t nonce, uint32_t *hashOutput) override;
     void set(const Job &job, uint8_t *blob) override;
 
 protected:
@@ -64,7 +63,6 @@ protected:
     Blake2bHashRegistersKernel *m_blake2b_hash_registers_64       = nullptr;
     Blake2bInitialHashKernel *m_blake2b_initial_hash              = nullptr;
     Blake2bInitialHashDoubleKernel *m_blake2b_initial_hash_double = nullptr;
-    Blake2bInitialHashBigKernel* m_blake2b_initial_hash_big       = nullptr;
     Buffer m_seed;
     cl_mem m_dataset                                              = nullptr;
     cl_mem m_entropy                                              = nullptr;

src/base/net/stratum/Client.cpp

@@ -363,7 +363,7 @@ bool xmrig::Client::parseJob(const rapidjson::Value &params, int *code)
 
     Job job(has<EXT_NICEHASH>(), m_pool.algorithm(), m_rpcId);
 
-    if (!job.setId(Json::getString(params, "job_id"))) {
+    if (!job.setId(params["job_id"].GetString())) {
         *code = 3;
         return false;
     }
@@ -400,7 +400,7 @@ bool xmrig::Client::parseJob(const rapidjson::Value &params, int *code)
         }
     }
 
-    if (!job.setTarget(Json::getString(params, "target"))) {
+    if (!job.setTarget(params["target"].GetString())) {
         *code = 5;
         return false;
     }

src/base/net/stratum/Job.cpp

@@ -110,12 +110,16 @@ bool xmrig::Job::setSeedHash(const char *hash)
 
 bool xmrig::Job::setTarget(const char *target)
 {
-    static auto parse = [](const char *target, size_t size, const Algorithm &algorithm) -> uint64_t {
+    static auto parse = [](const char *target, const Algorithm &algorithm) -> uint64_t {
+        if (!target) {
+            return 0;
+        }
+
         if (algorithm == Algorithm::RX_YADA) {
             return strtoull(target, nullptr, 16);
         }
 
-        const auto raw = Cvt::fromHex(target, size);
+        const auto raw = Cvt::fromHex(target, strlen(target));
 
         switch (raw.size()) {
         case 4:
@@ -131,48 +135,23 @@ bool xmrig::Job::setTarget(const char *target)
         return 0;
     };
 
-    const size_t size = target ? strlen(target) : 0;
+    if ((m_target = parse(target, algorithm())) == 0) {
-
-    if (size < 4 || (m_target = parse(target, size, algorithm())) == 0) {
         return false;
     }
 
     m_diff = toDiff(m_target);
 
 #   ifdef XMRIG_PROXY_PROJECT
-    if (size >= sizeof(m_rawTarget)) {
+    assert(sizeof(m_rawTarget) > (size * 2));
-        return false;
-    }
 
     memset(m_rawTarget, 0, sizeof(m_rawTarget));
-    memcpy(m_rawTarget, target, size);
+    memcpy(m_rawTarget, target, std::min(size * 2, sizeof(m_rawTarget)));
 #   endif
 
     return true;
 }
 
 
-size_t xmrig::Job::nonceOffset() const
-{
-    switch (algorithm().family()) {
-    case Algorithm::KAWPOW:
-        return 32;
-
-    case Algorithm::GHOSTRIDER:
-        return 76;
-
-    default:
-        break;
-    }
-
-    if (algorithm() == Algorithm::RX_YADA) {
-        return 147;
-    }
-
-    return 39;
-}
-
-
 void xmrig::Job::setDiff(uint64_t diff)
 {
     m_diff   = diff;
@@ -203,6 +182,26 @@ void xmrig::Job::setSigKey(const char *sig_key)
 }
 
 
+int32_t xmrig::Job::nonceOffset() const
+{
+    switch (algorithm().family()) {
+    case Algorithm::KAWPOW:
+        return 32;
+
+    case Algorithm::GHOSTRIDER:
+        return 76;
+
+    default:
+        break;
+    }
+
+    if (algorithm() == Algorithm::RX_YADA) {
+        return 147;
+    }
+
+    return 39;
+}
+
 uint32_t xmrig::Job::getNumTransactions() const
 {
     if (!(m_algorithm.isCN() || m_algorithm.family() == Algorithm::RANDOM_X)) {

src/base/net/stratum/Job.h

@@ -61,7 +61,6 @@ public:
     bool setBlob(const char *blob);
     bool setSeedHash(const char *hash);
     bool setTarget(const char *target);
-    size_t nonceOffset() const;
     void setDiff(uint64_t diff);
     void setSigKey(const char *sig_key);
 
@@ -76,6 +75,7 @@ public:
     inline const String &poolWallet() const             { return m_poolWallet; }
     inline const uint32_t *nonce() const                { return reinterpret_cast<const uint32_t*>(m_blob + nonceOffset()); }
     inline const uint8_t *blob() const                  { return m_blob; }
+    int32_t nonceOffset() const;
     inline size_t nonceSize() const                     { return (algorithm().family() == Algorithm::KAWPOW) ?  8 :  4; }
     inline size_t size() const                          { return m_size; }
     inline uint32_t *nonce()                            { return reinterpret_cast<uint32_t*>(m_blob + nonceOffset()); }

src/core/config/usage.h

@@ -94,7 +94,13 @@ static inline const std::string &usage()
 #   ifdef XMRIG_ALGO_RANDOMX
     u += "      --huge-pages-jit          enable huge pages support for RandomX JIT code\n";
 #   endif
+#   ifdef XMRIG_FEATURE_ASM
+#   ifdef XMRIG_FEATURE_ASM_AMD
     u += "      --asm=ASM                 ASM optimizations, possible values: auto, none, intel, ryzen, bulldozer\n";
+#   else
+    u += "      --asm=ASM                 ASM optimizations, possible values: auto, none, intel\n";
+#   endif
+#   endif
 
 #   if defined(__x86_64__) || defined(_M_AMD64)
     u += "      --argon2-impl=IMPL        argon2 implementation: x86_64, SSE2, SSSE3, XOP, AVX2, AVX-512F\n";

src/crypto/cn/CnHash.cpp

@@ -55,6 +55,7 @@ bool cn_vaes_enabled = false;
 
 
 #ifdef XMRIG_FEATURE_ASM
+#ifdef XMRIG_FEATURE_ASM_AMD
 #   define ADD_FN_ASM(algo) do {                                                                                    \
         m_map[algo]->data[AV_SINGLE][Assembly::INTEL]     = cryptonight_single_hash_asm<algo, Assembly::INTEL>;     \
         m_map[algo]->data[AV_SINGLE][Assembly::RYZEN]     = cryptonight_single_hash_asm<algo, Assembly::RYZEN>;     \
@@ -63,34 +64,50 @@ bool cn_vaes_enabled = false;
         m_map[algo]->data[AV_DOUBLE][Assembly::RYZEN]     = cryptonight_double_hash_asm<algo, Assembly::RYZEN>;     \
         m_map[algo]->data[AV_DOUBLE][Assembly::BULLDOZER] = cryptonight_double_hash_asm<algo, Assembly::BULLDOZER>; \
     } while (0)
+#else
+#   define ADD_FN_ASM(algo) do {                                                                                    \
+        m_map[algo]->data[AV_SINGLE][Assembly::INTEL]     = cryptonight_single_hash_asm<algo, Assembly::INTEL>;     \
+        m_map[algo]->data[AV_DOUBLE][Assembly::INTEL]     = cryptonight_double_hash_asm<algo, Assembly::INTEL>;     \
+    } while (0)
+#endif
 
 
 namespace xmrig {
 
 
 cn_mainloop_fun        cn_half_mainloop_ivybridge_asm             = nullptr;
+#ifdef XMRIG_FEATURE_ASM_AMD
 cn_mainloop_fun        cn_half_mainloop_ryzen_asm                 = nullptr;
 cn_mainloop_fun        cn_half_mainloop_bulldozer_asm             = nullptr;
+#endif
 cn_mainloop_fun        cn_half_double_mainloop_sandybridge_asm    = nullptr;
 
 cn_mainloop_fun        cn_trtl_mainloop_ivybridge_asm             = nullptr;
+#ifdef XMRIG_FEATURE_ASM_AMD
 cn_mainloop_fun        cn_trtl_mainloop_ryzen_asm                 = nullptr;
 cn_mainloop_fun        cn_trtl_mainloop_bulldozer_asm             = nullptr;
+#endif
 cn_mainloop_fun        cn_trtl_double_mainloop_sandybridge_asm    = nullptr;
 
 cn_mainloop_fun        cn_tlo_mainloop_ivybridge_asm              = nullptr;
+#ifdef XMRIG_FEATURE_ASM_AMD
 cn_mainloop_fun        cn_tlo_mainloop_ryzen_asm                  = nullptr;
 cn_mainloop_fun        cn_tlo_mainloop_bulldozer_asm              = nullptr;
+#endif
 cn_mainloop_fun        cn_tlo_double_mainloop_sandybridge_asm     = nullptr;
 
 cn_mainloop_fun        cn_zls_mainloop_ivybridge_asm              = nullptr;
+#ifdef XMRIG_FEATURE_ASM_AMD
 cn_mainloop_fun        cn_zls_mainloop_ryzen_asm                  = nullptr;
 cn_mainloop_fun        cn_zls_mainloop_bulldozer_asm              = nullptr;
+#endif
 cn_mainloop_fun        cn_zls_double_mainloop_sandybridge_asm     = nullptr;
 
 cn_mainloop_fun        cn_double_mainloop_ivybridge_asm           = nullptr;
+#ifdef XMRIG_FEATURE_ASM_AMD
 cn_mainloop_fun        cn_double_mainloop_ryzen_asm               = nullptr;
 cn_mainloop_fun        cn_double_mainloop_bulldozer_asm           = nullptr;
+#endif
 cn_mainloop_fun        cn_double_double_mainloop_sandybridge_asm  = nullptr;
 
 cn_mainloop_fun        cn_upx2_mainloop_asm                       = nullptr;
@@ -160,31 +177,41 @@ static void patchAsmVariants()
     auto base = static_cast<uint8_t *>(VirtualMemory::allocateExecutableMemory(allocation_size, false));
 
     cn_half_mainloop_ivybridge_asm              = reinterpret_cast<cn_mainloop_fun>         (base + 0x0000);
+#   ifdef XMRIG_FEATURE_ASM_AMD
     cn_half_mainloop_ryzen_asm                  = reinterpret_cast<cn_mainloop_fun>         (base + 0x1000);
     cn_half_mainloop_bulldozer_asm              = reinterpret_cast<cn_mainloop_fun>         (base + 0x2000);
+#   endif
     cn_half_double_mainloop_sandybridge_asm     = reinterpret_cast<cn_mainloop_fun>         (base + 0x3000);
 
 #   ifdef XMRIG_ALGO_CN_PICO
     cn_trtl_mainloop_ivybridge_asm              = reinterpret_cast<cn_mainloop_fun>         (base + 0x4000);
+#   ifdef XMRIG_FEATURE_ASM_AMD
     cn_trtl_mainloop_ryzen_asm                  = reinterpret_cast<cn_mainloop_fun>         (base + 0x5000);
     cn_trtl_mainloop_bulldozer_asm              = reinterpret_cast<cn_mainloop_fun>         (base + 0x6000);
+#   endif
     cn_trtl_double_mainloop_sandybridge_asm     = reinterpret_cast<cn_mainloop_fun>         (base + 0x7000);
 #   endif
 
     cn_zls_mainloop_ivybridge_asm               = reinterpret_cast<cn_mainloop_fun>         (base + 0x8000);
+#   ifdef XMRIG_FEATURE_ASM_AMD
     cn_zls_mainloop_ryzen_asm                   = reinterpret_cast<cn_mainloop_fun>         (base + 0x9000);
     cn_zls_mainloop_bulldozer_asm               = reinterpret_cast<cn_mainloop_fun>         (base + 0xA000);
+#   endif
     cn_zls_double_mainloop_sandybridge_asm      = reinterpret_cast<cn_mainloop_fun>         (base + 0xB000);
 
     cn_double_mainloop_ivybridge_asm            = reinterpret_cast<cn_mainloop_fun>         (base + 0xC000);
+#   ifdef XMRIG_FEATURE_ASM_AMD
     cn_double_mainloop_ryzen_asm                = reinterpret_cast<cn_mainloop_fun>         (base + 0xD000);
     cn_double_mainloop_bulldozer_asm            = reinterpret_cast<cn_mainloop_fun>         (base + 0xE000);
+#   endif
     cn_double_double_mainloop_sandybridge_asm   = reinterpret_cast<cn_mainloop_fun>         (base + 0xF000);
 
 #   ifdef XMRIG_ALGO_CN_PICO
     cn_tlo_mainloop_ivybridge_asm               = reinterpret_cast<cn_mainloop_fun>         (base + 0x10000);
+#   ifdef XMRIG_FEATURE_ASM_AMD
     cn_tlo_mainloop_ryzen_asm                   = reinterpret_cast<cn_mainloop_fun>         (base + 0x11000);
     cn_tlo_mainloop_bulldozer_asm               = reinterpret_cast<cn_mainloop_fun>         (base + 0x12000);
+#   endif
     cn_tlo_double_mainloop_sandybridge_asm      = reinterpret_cast<cn_mainloop_fun>         (base + 0x13000);
 #   endif
 
@@ -220,8 +247,10 @@ static void patchAsmVariants()
         constexpr uint32_t ITER = CnAlgo<Algorithm::CN_HALF>().iterations();
 
         patchCode(cn_half_mainloop_ivybridge_asm,            cnv2_mainloop_ivybridge_asm,           ITER);
+#       ifdef XMRIG_FEATURE_ASM_AMD
         patchCode(cn_half_mainloop_ryzen_asm,                cnv2_mainloop_ryzen_asm,               ITER);
         patchCode(cn_half_mainloop_bulldozer_asm,            cnv2_mainloop_bulldozer_asm,           ITER);
+#       endif
         patchCode(cn_half_double_mainloop_sandybridge_asm,   cnv2_double_mainloop_sandybridge_asm,  ITER);
     }
 
@@ -231,8 +260,10 @@ static void patchAsmVariants()
         constexpr uint32_t MASK = CnAlgo<Algorithm::CN_PICO_0>().mask();
 
         patchCode(cn_trtl_mainloop_ivybridge_asm,            cnv2_mainloop_ivybridge_asm,           ITER,   MASK);
+#       ifdef XMRIG_FEATURE_ASM_AMD
         patchCode(cn_trtl_mainloop_ryzen_asm,                cnv2_mainloop_ryzen_asm,               ITER,   MASK);
         patchCode(cn_trtl_mainloop_bulldozer_asm,            cnv2_mainloop_bulldozer_asm,           ITER,   MASK);
+#       endif
         patchCode(cn_trtl_double_mainloop_sandybridge_asm,   cnv2_double_mainloop_sandybridge_asm,  ITER,   MASK);
     }
 
@@ -241,8 +272,10 @@ static void patchAsmVariants()
         constexpr uint32_t MASK = CnAlgo<Algorithm::CN_PICO_TLO>().mask();
 
         patchCode(cn_tlo_mainloop_ivybridge_asm,             cnv2_mainloop_ivybridge_asm,           ITER,   MASK);
+#       ifdef XMRIG_FEATURE_ASM_AMD
         patchCode(cn_tlo_mainloop_ryzen_asm,                 cnv2_mainloop_ryzen_asm,               ITER,   MASK);
         patchCode(cn_tlo_mainloop_bulldozer_asm,             cnv2_mainloop_bulldozer_asm,           ITER,   MASK);
+#       endif
         patchCode(cn_tlo_double_mainloop_sandybridge_asm,    cnv2_double_mainloop_sandybridge_asm,  ITER,   MASK);
     }
 #   endif
@@ -251,8 +284,10 @@ static void patchAsmVariants()
         constexpr uint32_t ITER = CnAlgo<Algorithm::CN_ZLS>().iterations();
 
         patchCode(cn_zls_mainloop_ivybridge_asm,             cnv2_mainloop_ivybridge_asm,           ITER);
+#       ifdef XMRIG_FEATURE_ASM_AMD
         patchCode(cn_zls_mainloop_ryzen_asm,                 cnv2_mainloop_ryzen_asm,               ITER);
         patchCode(cn_zls_mainloop_bulldozer_asm,             cnv2_mainloop_bulldozer_asm,           ITER);
+#       endif
         patchCode(cn_zls_double_mainloop_sandybridge_asm,    cnv2_double_mainloop_sandybridge_asm,  ITER);
     }
 
@@ -260,8 +295,10 @@ static void patchAsmVariants()
         constexpr uint32_t ITER = CnAlgo<Algorithm::CN_DOUBLE>().iterations();
 
         patchCode(cn_double_mainloop_ivybridge_asm,          cnv2_mainloop_ivybridge_asm,           ITER);
+#       ifdef XMRIG_FEATURE_ASM_AMD
         patchCode(cn_double_mainloop_ryzen_asm,              cnv2_mainloop_ryzen_asm,               ITER);
         patchCode(cn_double_mainloop_bulldozer_asm,          cnv2_mainloop_bulldozer_asm,           ITER);
+#       endif
         patchCode(cn_double_double_mainloop_sandybridge_asm, cnv2_double_mainloop_sandybridge_asm,  ITER);
     }
 

src/crypto/cn/CryptoNight_x86.h

@@ -852,12 +852,16 @@ extern "C" void cnv1_single_mainloop_asm(cryptonight_ctx * *ctx);
 extern "C" void cnv1_double_mainloop_asm(cryptonight_ctx **ctx);
 extern "C" void cnv1_quad_mainloop_asm(cryptonight_ctx **ctx);
 extern "C" void cnv2_mainloop_ivybridge_asm(cryptonight_ctx **ctx);
+#ifdef XMRIG_FEATURE_ASM_AMD
 extern "C" void cnv2_mainloop_ryzen_asm(cryptonight_ctx **ctx);
 extern "C" void cnv2_mainloop_bulldozer_asm(cryptonight_ctx **ctx);
+#endif
 extern "C" void cnv2_double_mainloop_sandybridge_asm(cryptonight_ctx **ctx);
 extern "C" void cnv2_rwz_mainloop_asm(cryptonight_ctx **ctx);
 extern "C" void cnv2_rwz_double_mainloop_asm(cryptonight_ctx **ctx);
+#ifdef XMRIG_FEATURE_ASM_AMD
 extern "C" void cnv2_upx_double_mainloop_zen3_asm(cryptonight_ctx **ctx);
+#endif
 
 
 namespace xmrig {
@@ -867,28 +871,38 @@ typedef void (*cn_mainloop_fun)(cryptonight_ctx **ctx);
 
 
 extern cn_mainloop_fun cn_half_mainloop_ivybridge_asm;
+#ifdef XMRIG_FEATURE_ASM_AMD
 extern cn_mainloop_fun cn_half_mainloop_ryzen_asm;
 extern cn_mainloop_fun cn_half_mainloop_bulldozer_asm;
+#endif
 extern cn_mainloop_fun cn_half_double_mainloop_sandybridge_asm;
 
 extern cn_mainloop_fun cn_trtl_mainloop_ivybridge_asm;
+#ifdef XMRIG_FEATURE_ASM_AMD
 extern cn_mainloop_fun cn_trtl_mainloop_ryzen_asm;
 extern cn_mainloop_fun cn_trtl_mainloop_bulldozer_asm;
+#endif
 extern cn_mainloop_fun cn_trtl_double_mainloop_sandybridge_asm;
 
 extern cn_mainloop_fun cn_tlo_mainloop_ivybridge_asm;
+#ifdef XMRIG_FEATURE_ASM_AMD
 extern cn_mainloop_fun cn_tlo_mainloop_ryzen_asm;
 extern cn_mainloop_fun cn_tlo_mainloop_bulldozer_asm;
+#endif
 extern cn_mainloop_fun cn_tlo_double_mainloop_sandybridge_asm;
 
 extern cn_mainloop_fun cn_zls_mainloop_ivybridge_asm;
+#ifdef XMRIG_FEATURE_ASM_AMD
 extern cn_mainloop_fun cn_zls_mainloop_ryzen_asm;
 extern cn_mainloop_fun cn_zls_mainloop_bulldozer_asm;
+#endif
 extern cn_mainloop_fun cn_zls_double_mainloop_sandybridge_asm;
 
 extern cn_mainloop_fun cn_double_mainloop_ivybridge_asm;
+#ifdef XMRIG_FEATURE_ASM_AMD
 extern cn_mainloop_fun cn_double_mainloop_ryzen_asm;
 extern cn_mainloop_fun cn_double_mainloop_bulldozer_asm;
+#endif
 extern cn_mainloop_fun cn_double_double_mainloop_sandybridge_asm;
 
 extern cn_mainloop_fun cn_upx2_mainloop_asm;
@@ -964,46 +978,54 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
         if (ASM == Assembly::INTEL) {
             cnv2_mainloop_ivybridge_asm(ctx);
         }
+#       ifdef XMRIG_FEATURE_ASM_AMD
         else if (ASM == Assembly::RYZEN) {
             cnv2_mainloop_ryzen_asm(ctx);
         }
         else {
             cnv2_mainloop_bulldozer_asm(ctx);
         }
+#       endif
     }
     else if (ALGO == Algorithm::CN_HALF) {
         if (ASM == Assembly::INTEL) {
             cn_half_mainloop_ivybridge_asm(ctx);
         }
+#       ifdef XMRIG_FEATURE_ASM_AMD
         else if (ASM == Assembly::RYZEN) {
             cn_half_mainloop_ryzen_asm(ctx);
         }
         else {
             cn_half_mainloop_bulldozer_asm(ctx);
         }
+#       endif
     }
 #   ifdef XMRIG_ALGO_CN_PICO
     else if (ALGO == Algorithm::CN_PICO_0) {
         if (ASM == Assembly::INTEL) {
             cn_trtl_mainloop_ivybridge_asm(ctx);
         }
+#       ifdef XMRIG_FEATURE_ASM_AMD
         else if (ASM == Assembly::RYZEN) {
             cn_trtl_mainloop_ryzen_asm(ctx);
         }
         else {
             cn_trtl_mainloop_bulldozer_asm(ctx);
         }
+#       endif
     }
     else if (ALGO == Algorithm::CN_PICO_TLO) {
         if (ASM == Assembly::INTEL) {
             cn_tlo_mainloop_ivybridge_asm(ctx);
         }
+#       ifdef XMRIG_FEATURE_ASM_AMD
         else if (ASM == Assembly::RYZEN) {
             cn_tlo_mainloop_ryzen_asm(ctx);
         }
         else {
             cn_tlo_mainloop_bulldozer_asm(ctx);
         }
+#       endif
     }
 #   endif
     else if (ALGO == Algorithm::CN_RWZ) {
@@ -1013,23 +1035,27 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
         if (ASM == Assembly::INTEL) {
             cn_zls_mainloop_ivybridge_asm(ctx);
         }
+#       ifdef XMRIG_FEATURE_ASM_AMD
         else if (ASM == Assembly::RYZEN) {
             cn_zls_mainloop_ryzen_asm(ctx);
         }
         else {
             cn_zls_mainloop_bulldozer_asm(ctx);
         }
+#       endif
     }
     else if (ALGO == Algorithm::CN_DOUBLE) {
         if (ASM == Assembly::INTEL) {
             cn_double_mainloop_ivybridge_asm(ctx);
         }
+#       ifdef XMRIG_FEATURE_ASM_AMD
         else if (ASM == Assembly::RYZEN) {
             cn_double_mainloop_ryzen_asm(ctx);
         }
         else {
             cn_double_mainloop_bulldozer_asm(ctx);
         }
+#       endif
     }
 #   ifdef XMRIG_ALGO_CN_FEMTO
     else if (ALGO == Algorithm::CN_UPX2) {
@@ -1094,12 +1120,16 @@ inline void cryptonight_double_hash_asm(const uint8_t *__restrict__ input, size_
 #   endif
 #   ifdef XMRIG_ALGO_CN_FEMTO
     else if (ALGO == Algorithm::CN_UPX2) {
+#       ifdef XMRIG_FEATURE_ASM_AMD
         if (Cpu::info()->arch() == ICpuInfo::ARCH_ZEN3) {
             cnv2_upx_double_mainloop_zen3_asm(ctx);
         }
         else {
             cn_upx2_double_mainloop_asm(ctx);
         }
+#       else
+        cn_upx2_double_mainloop_asm(ctx);
+#       endif
     }
 #   endif
     else if (ALGO == Algorithm::CN_RWZ) {

src/crypto/cn/asm/cn_main_loop.S

@@ -15,12 +15,16 @@
 .global FN_PREFIX(cnv1_double_mainloop_asm)
 .global FN_PREFIX(cnv1_quad_mainloop_asm)
 .global FN_PREFIX(cnv2_mainloop_ivybridge_asm)
+#ifdef XMRIG_FEATURE_ASM_AMD
 .global FN_PREFIX(cnv2_mainloop_ryzen_asm)
 .global FN_PREFIX(cnv2_mainloop_bulldozer_asm)
+#endif
 .global FN_PREFIX(cnv2_double_mainloop_sandybridge_asm)
 .global FN_PREFIX(cnv2_rwz_mainloop_asm)
 .global FN_PREFIX(cnv2_rwz_double_mainloop_asm)
+#ifdef XMRIG_FEATURE_ASM_AMD
 .global FN_PREFIX(cnv2_upx_double_mainloop_zen3_asm)
+#endif
 
 ALIGN(64)
 FN_PREFIX(cnv1_single_mainloop_asm):
@@ -58,6 +62,7 @@ FN_PREFIX(cnv2_mainloop_ivybridge_asm):
 	ret 0
 	mov eax, 3735929054
 
+#ifdef XMRIG_FEATURE_ASM_AMD
 ALIGN(64)
 FN_PREFIX(cnv2_mainloop_ryzen_asm):
 	sub rsp, 48
@@ -75,6 +80,7 @@ FN_PREFIX(cnv2_mainloop_bulldozer_asm):
 	add rsp, 48
 	ret 0
 	mov eax, 3735929054
+#endif
 
 ALIGN(64)
 FN_PREFIX(cnv2_double_mainloop_sandybridge_asm):
@@ -103,6 +109,7 @@ FN_PREFIX(cnv2_rwz_double_mainloop_asm):
 	ret 0
 	mov eax, 3735929054
 
+#ifdef XMRIG_FEATURE_ASM_AMD
 ALIGN(64)
 FN_PREFIX(cnv2_upx_double_mainloop_zen3_asm):
 	sub rsp, 48
@@ -111,6 +118,7 @@ FN_PREFIX(cnv2_upx_double_mainloop_zen3_asm):
 	add rsp, 48
 	ret 0
 	mov eax, 3735929054
+#endif
 
 #if defined(__linux__) && defined(__ELF__)
 .section .note.GNU-stack,"",%progbits

src/crypto/cn/asm/win64/cn_main_loop.S

@@ -5,12 +5,16 @@
 .global cnv1_double_mainloop_asm
 .global cnv1_quad_mainloop_asm
 .global cnv2_mainloop_ivybridge_asm
+#ifdef XMRIG_FEATURE_ASM_AMD
 .global cnv2_mainloop_ryzen_asm
 .global cnv2_mainloop_bulldozer_asm
+#endif
 .global cnv2_double_mainloop_sandybridge_asm
 .global cnv2_rwz_mainloop_asm
 .global cnv2_rwz_double_mainloop_asm
+#ifdef XMRIG_FEATURE_ASM_AMD
 .global cnv2_upx_double_mainloop_zen3_asm
+#endif
 
 ALIGN(64)
 cnv1_single_mainloop_asm:
@@ -36,6 +40,7 @@ cnv2_mainloop_ivybridge_asm:
 	ret 0
 	mov eax, 3735929054
 
+#ifdef XMRIG_FEATURE_ASM_AMD
 ALIGN(64)
 cnv2_mainloop_ryzen_asm:
 	#include "../cn2/cnv2_main_loop_ryzen.inc"
@@ -47,6 +52,7 @@ cnv2_mainloop_bulldozer_asm:
 	#include "../cn2/cnv2_main_loop_bulldozer.inc"
 	ret 0
 	mov eax, 3735929054
+#endif
 
 ALIGN(64)
 cnv2_double_mainloop_sandybridge_asm:
@@ -66,8 +72,10 @@ cnv2_rwz_double_mainloop_asm:
 	ret 0
 	mov eax, 3735929054
 
+#ifdef XMRIG_FEATURE_ASM_AMD
 ALIGN(64)
 cnv2_upx_double_mainloop_zen3_asm:
 	#include "cn2/cnv2_upx_double_mainloop_zen3.inc"
 	ret 0
 	mov eax, 3735929054
+#endif

src/crypto/rx/RxVm.cpp

@@ -41,10 +41,12 @@ randomx_vm *xmrig::RxVm::create(RxDataset *dataset, uint8_t *scratchpad, bool so
         flags |= RANDOMX_FLAG_JIT;
     }
 
+#   ifdef XMRIG_FEATURE_ASM_AMD
     const auto asmId = assembly == Assembly::AUTO ? Cpu::info()->assembly() : assembly.id();
     if ((asmId == Assembly::RYZEN) || (asmId == Assembly::BULLDOZER)) {
         flags |= RANDOMX_FLAG_AMD;
     }
+#   endif
 
     return randomx_create_vm(static_cast<randomx_flags>(flags), !dataset->get() ? dataset->cache()->get() : nullptr, dataset->get(), scratchpad, node);
 }

src/version.h

@@ -22,15 +22,15 @@
 #define APP_ID        "xmrig"
 #define APP_NAME      "XMRig"
 #define APP_DESC      "XMRig miner"
-#define APP_VERSION   "6.22.0"
+#define APP_VERSION   "6.21.4-dev"
 #define APP_DOMAIN    "xmrig.com"
 #define APP_SITE      "www.xmrig.com"
 #define APP_COPYRIGHT "Copyright (C) 2016-2024 xmrig.com"
 #define APP_KIND      "miner"
 
 #define APP_VER_MAJOR  6
-#define APP_VER_MINOR  22
+#define APP_VER_MINOR  21
-#define APP_VER_PATCH  0
+#define APP_VER_PATCH  4
 
 #ifdef _MSC_VER
 #   if (_MSC_VER >= 1930)