Merge ab5be0b773 into 16ecb8f085

CUDA backend: update RandomX dataset when it changes

CHANGELOG.md

@@ -1,3 +1,22 @@
+# v6.22.2
+- [#3569](https://github.com/xmrig/xmrig/pull/3569) Fixed corrupted API output in some rare conditions.
+- [#3571](https://github.com/xmrig/xmrig/pull/3571) Fixed number of threads on the new Intel Core Ultra CPUs.
+
+# v6.22.1
+- [#3531](https://github.com/xmrig/xmrig/pull/3531) Always reset nonce on RandomX dataset change.
+- [#3534](https://github.com/xmrig/xmrig/pull/3534) Fixed threads auto-config on Zen5.
+- [#3535](https://github.com/xmrig/xmrig/pull/3535) RandomX: tweaks for Zen5.
+- [#3539](https://github.com/xmrig/xmrig/pull/3539) Added Zen5 to `randomx_boost.sh`.
+- [#3540](https://github.com/xmrig/xmrig/pull/3540) Detect AMD engineering samples in `randomx_boost.sh`.
+
+# 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.
 

doc/ALGORITHMS.md

@@ -13,7 +13,6 @@ Option `coin` useful for pools without [algorithm negotiation](https://xmrig.com
 | Name | Memory | Version | Description | Notes |
 |------|--------|---------|-------------|-------|
 | `kawpow` | - | 6.0.0+ | KawPow (Ravencoin) | GPU only |
-| `rx/keva` | 1 MB | 5.9.0+ | RandomKEVA (RandomX variant for Keva). |  |
 | `astrobwt` | 20 MB | 5.8.0+ | AstroBWT (Dero). |  |
 | `cn-pico/tlo` | 256 KB | 5.5.0+ | CryptoNight-Pico (Talleo). |  |
 | `rx/sfx` | 2 MB | 5.4.0+ | RandomSFX (RandomX variant for Safex). |  |

scripts/build.hwloc.sh

@@ -1,8 +1,8 @@
 #!/bin/sh -e
 
 HWLOC_VERSION_MAJOR="2"
-HWLOC_VERSION_MINOR="10"
-HWLOC_VERSION_PATCH="0"
+HWLOC_VERSION_MINOR="11"
+HWLOC_VERSION_PATCH="2"
 
 HWLOC_VERSION="${HWLOC_VERSION_MAJOR}.${HWLOC_VERSION_MINOR}.${HWLOC_VERSION_PATCH}"
 

scripts/build.openssl.sh

@@ -1,6 +1,6 @@
 #!/bin/sh -e
 
-OPENSSL_VERSION="1.1.1s"
+OPENSSL_VERSION="1.1.1u"
 
 mkdir -p deps
 mkdir -p deps/include
@@ -8,7 +8,7 @@ mkdir -p deps/lib
 
 mkdir -p build && cd build
 
-wget https://www.openssl.org/source/openssl-${OPENSSL_VERSION}.tar.gz -O openssl-${OPENSSL_VERSION}.tar.gz
+wget https://openssl.org/source/old/1.1.1/openssl-${OPENSSL_VERSION}.tar.gz -O openssl-${OPENSSL_VERSION}.tar.gz
 tar -xzf openssl-${OPENSSL_VERSION}.tar.gz
 
 cd openssl-${OPENSSL_VERSION}

scripts/build.openssl3.sh

@@ -1,6 +1,6 @@
 #!/bin/sh -e
 
-OPENSSL_VERSION="3.0.13"
+OPENSSL_VERSION="3.0.15"
 
 mkdir -p deps
 mkdir -p deps/include
@@ -8,7 +8,7 @@ mkdir -p deps/lib
 
 mkdir -p build && cd build
 
-wget https://www.openssl.org/source/openssl-${OPENSSL_VERSION}.tar.gz -O openssl-${OPENSSL_VERSION}.tar.gz
+wget https://github.com/openssl/openssl/releases/download/openssl-${OPENSSL_VERSION}/openssl-${OPENSSL_VERSION}.tar.gz -O openssl-${OPENSSL_VERSION}.tar.gz
 tar -xzf openssl-${OPENSSL_VERSION}.tar.gz
 
 cd openssl-${OPENSSL_VERSION}

scripts/build.uv.sh

@@ -1,6 +1,6 @@
 #!/bin/sh -e
 
-UV_VERSION="1.48.0"
+UV_VERSION="1.49.2"
 
 mkdir -p deps
 mkdir -p deps/include

scripts/generate_cl.js

@@ -50,7 +50,6 @@ function rx()
         'randomx_constants_monero.h',
         'randomx_constants_wow.h',
         'randomx_constants_arqma.h',
-        'randomx_constants_keva.h',
         'randomx_constants_graft.h',
         'aes.cl',
         'blake2b.cl',

scripts/randomx_boost.sh

@@ -8,7 +8,7 @@ else
 	modprobe msr allow_writes=on
 fi
 
-if grep -E 'AMD Ryzen|AMD EPYC' /proc/cpuinfo > /dev/null;
+if grep -E 'AMD Ryzen|AMD EPYC|AuthenticAMD' /proc/cpuinfo > /dev/null;
 	then
 	if grep "cpu family[[:space:]]\{1,\}:[[:space:]]25" /proc/cpuinfo > /dev/null;
 		then
@@ -28,6 +28,14 @@ if grep -E 'AMD Ryzen|AMD EPYC' /proc/cpuinfo > /dev/null;
 					wrmsr -a 0xc001102b 0x2000cc10
 					echo "MSR register values for Zen3 applied"
 				fi
+		elif grep "cpu family[[:space:]]\{1,\}:[[:space:]]26" /proc/cpuinfo > /dev/null;
+			then
+				echo "Detected Zen5 CPU"
+				wrmsr -a 0xc0011020 0x4400000000000
+				wrmsr -a 0xc0011021 0x4000000000040
+				wrmsr -a 0xc0011022 0x8680000401570000
+				wrmsr -a 0xc001102b 0x2040cc10
+				echo "MSR register values for Zen5 applied"
 		else
 			echo "Detected Zen1/Zen2 CPU"
 			wrmsr -a 0xc0011020 0

src/3rdparty/hwloc/NEWS

@@ -1,5 +1,5 @@
 Copyright © 2009 CNRS
-Copyright © 2009-2023 Inria.  All rights reserved.
+Copyright © 2009-2024 Inria.  All rights reserved.
 Copyright © 2009-2013 Université Bordeaux
 Copyright © 2009-2011 Cisco Systems, Inc.  All rights reserved.
 Copyright © 2020 Hewlett Packard Enterprise.  All rights reserved.
@@ -17,6 +17,71 @@ bug fixes (and other actions) for each version of hwloc since version
 0.9.
 
 
+Version 2.11.2
+--------------
+* Add missing CPU info attrs on aarch64 on Linux.
+* Use ACPI CPPC on Linux to get better information about cpukinds,
+  at least on AMD CPUs.
+* Fix crash when manipulating cpukinds after topology
+  duplication, thanks to Hadrien Grasland for the report.
+* Fix missing input target checks in memattr functions,
+  thanks to Hadrien Grasland for the report.
+* Fix a memory leak when ignoring NUMA distances on FreeBSD.
+* Fix build failure on old Linux distributions without accessat().
+* Fix non-Windows importing of XML topologies and CPUID dumps exported
+  on Windows.
+* hwloc-calc --cpuset-output-format systemd-dbus-api now allows
+  to generate AllowedCPUs information for systemd slices.
+  See the hwloc-calc manpage for examples. Thanks to Pierre Neyron.
+* Some fixes in manpage EXAMPLES and split them into subsections.
+
+
+Version 2.11.1
+--------------
+* Fix bash completions, thanks Tavis Rudd.
+
+
+Version 2.11.0
+--------------
+* API
+  + Add HWLOC_MEMBIND_WEIGHTED_INTERLEAVE memory binding policy on
+    Linux 6.9+. Thanks to Honggyu Kim for the patch.
+    - weighted_interleave_membind is added to membind support bits.
+    - The "weighted" policy is added to the hwloc-bind tool.
+  + Add hwloc_obj_set_subtype(). Thanks to Hadrien Grasland for the report.
+* GPU support
+  + Don't hide the GPU NUMA node on NVIDIA Grace Hopper.
+  + Get Intel GPU OpenCL device locality.
+  + Add bandwidths between subdevices in the LevelZero XeLinkBandwidth
+    matrix.
+  + Fix PCI Gen4+ link speed of NVIDIA GPU obtained from NVML,
+    thanks to Akram Sbaih for the report.
+* Windows support
+  + Fix Windows support when UNICODE is enabled, several hwloc features
+    were missing, thanks to Martin for the report.
+  + Fix the enabling of CUDA in Windows CMake build,
+    Thanks to Moritz Kreutzer for the patch.
+  + Fix CUDA/OpenCL test source path in Windows CMake.
+* Tools
+  + Option --best-memattr may now return multiple nodes. Additional
+    configuration flags may be given to tweak its behavior.
+  + hwloc-info has a new --get-attr option to get a single attribute.
+  + hwloc-info now supports "levels", "support" and "topology"
+    special keywords for backward compatibility for hwloc 3.0.
+  + The --taskset command-line option is superseded by the new
+    --cpuset-output-format which also allows to export as list.
+  + hwloc-calc may now import bitmasks described as a list of bits
+    with the new "--cpuset-input-format list".
+* Misc
+  + The MemoryTiersNr info attribute in the root object now says how many
+    memory tiers were built. Thanks to Antoine Morvan for the report.
+  + Fix the management of infinite cpusets in the bitmap printf/sscanf
+    API as well as in command-line tools.
+  + Add section "Compiling software on top of hwloc's C API" in the
+    documentation with examples for GNU Make and CMake,
+    thanks to Florent Pruvost for the help.
+
+
 Version 2.10.0
 --------------
 * Heterogeneous Memory core improvements

src/3rdparty/hwloc/README

@@ -418,14 +418,8 @@ return 0;
 }
 
 hwloc provides a pkg-config executable to obtain relevant compiler and linker
-flags. For example, it can be used thusly to compile applications that utilize
-the hwloc library (assuming GNU Make):
-
-CFLAGS += $(shell pkg-config --cflags hwloc)
-LDLIBS += $(shell pkg-config --libs hwloc)
-
-hwloc-hello: hwloc-hello.c
-     $(CC) hwloc-hello.c $(CFLAGS) -o hwloc-hello $(LDLIBS)
+flags. See Compiling software on top of hwloc's C API for details on building
+program on top of hwloc's API using GNU Make or CMake.
 
 On a machine 2 processor packages -- each package of which has two processing
 cores -- the output from running hwloc-hello could be something like the

src/3rdparty/hwloc/VERSION

@@ -8,8 +8,8 @@
 # Please update HWLOC_VERSION* in contrib/windows/hwloc_config.h too.
 
 major=2
-minor=10
-release=0
+minor=11
+release=2
 
 # greek is used for alpha or beta release tags.  If it is non-empty,
 # it will be appended to the version number.  It does not have to be
@@ -22,7 +22,7 @@ greek=
 
 # The date when this release was created
 
-date="Dec 04, 2023"
+date="Sep 26, 2024"
 
 # If snapshot=1, then use the value from snapshot_version as the
 # entire hwloc version (i.e., ignore major, minor, release, and
@@ -41,6 +41,6 @@ snapshot_version=${major}.${minor}.${release}${greek}-git
 # 2. Version numbers are described in the Libtool current:revision:age
 # format.
 
-libhwloc_so_version=22:0:7
+libhwloc_so_version=23:1:8
 
 # Please also update the <TargetName> lines in contrib/windows/libhwloc.vcxproj

src/3rdparty/hwloc/include/hwloc/autogen/config.h

@@ -1,6 +1,6 @@
 /*
  * Copyright © 2009 CNRS
- * Copyright © 2009-2023 Inria.  All rights reserved.
+ * Copyright © 2009-2024 Inria.  All rights reserved.
  * Copyright © 2009-2012 Université Bordeaux
  * Copyright © 2009-2011 Cisco Systems, Inc.  All rights reserved.
  * See COPYING in top-level directory.
@@ -11,10 +11,10 @@
 #ifndef HWLOC_CONFIG_H
 #define HWLOC_CONFIG_H
 
-#define HWLOC_VERSION "2.10.0"
+#define HWLOC_VERSION "2.11.2"
 #define HWLOC_VERSION_MAJOR 2
-#define HWLOC_VERSION_MINOR 10
-#define HWLOC_VERSION_RELEASE 0
+#define HWLOC_VERSION_MINOR 11
+#define HWLOC_VERSION_RELEASE 2
 #define HWLOC_VERSION_GREEK ""
 
 #define __hwloc_restrict

src/3rdparty/hwloc/include/hwloc/distances.h

@@ -1,5 +1,5 @@
 /*
- * Copyright © 2010-2023 Inria.  All rights reserved.
+ * Copyright © 2010-2024 Inria.  All rights reserved.
  * See COPYING in top-level directory.
  */
 
@@ -28,18 +28,18 @@ extern "C" {
 
 /** \brief Matrix of distances between a set of objects.
  *
- * This matrix often contains latencies between NUMA nodes
+ * The most common matrix contains latencies between NUMA nodes
  * (as reported in the System Locality Distance Information Table (SLIT)
  * in the ACPI specification), which may or may not be physically accurate.
  * It corresponds to the latency for accessing the memory of one node
  * from a core in another node.
- * The corresponding kind is ::HWLOC_DISTANCES_KIND_FROM_OS | ::HWLOC_DISTANCES_KIND_FROM_USER.
+ * The corresponding kind is ::HWLOC_DISTANCES_KIND_MEANS_LATENCY | ::HWLOC_DISTANCES_KIND_FROM_USER.
  * The name of this distances structure is "NUMALatency".
- * Others distance structures include and "XGMIBandwidth", "XGMIHops",
- * "XeLinkBandwidth" and "NVLinkBandwidth".
  *
  * The matrix may also contain bandwidths between random sets of objects,
  * possibly provided by the user, as specified in the \p kind attribute.
+ * Others common distance structures include and "XGMIBandwidth", "XGMIHops",
+ * "XeLinkBandwidth" and "NVLinkBandwidth".
  *
  * Pointers \p objs and \p values should not be replaced, reallocated, freed, etc.
  * However callers are allowed to modify \p kind as well as the contents
@@ -70,11 +70,10 @@ struct hwloc_distances_s {
  * The \p kind attribute of struct hwloc_distances_s is a OR'ed set
  * of kinds.
  *
- * A kind of format HWLOC_DISTANCES_KIND_FROM_* specifies where the
- * distance information comes from, if known.
- *
- * A kind of format HWLOC_DISTANCES_KIND_MEANS_* specifies whether
- * values are latencies or bandwidths, if applicable.
+ * Each distance matrix may have only one kind among HWLOC_DISTANCES_KIND_FROM_*
+ * specifying where distance information comes from,
+ * and one kind among HWLOC_DISTANCES_KIND_MEANS_* specifying
+ * whether values are latencies or bandwidths.
  */
 enum hwloc_distances_kind_e {
   /** \brief These distances were obtained from the operating system or hardware.
@@ -357,6 +356,8 @@ typedef void * hwloc_distances_add_handle_t;
  * Otherwise, it will be copied internally and may later be freed by the caller.
  *
  * \p kind specifies the kind of distance as a OR'ed set of ::hwloc_distances_kind_e.
+ * Only one kind of meaning and one kind of provenance may be given if appropriate
+ * (e.g. ::HWLOC_DISTANCES_KIND_MEANS_BANDWIDTH and ::HWLOC_DISTANCES_KIND_FROM_USER).
  * Kind ::HWLOC_DISTANCES_KIND_HETEROGENEOUS_TYPES will be automatically set
  * according to objects having different types in hwloc_distances_add_values().
  *
@@ -403,7 +404,8 @@ HWLOC_DECLSPEC int hwloc_distances_add_values(hwloc_topology_t topology,
 /** \brief Flags for adding a new distances to a topology. */
 enum hwloc_distances_add_flag_e {
   /** \brief Try to group objects based on the newly provided distance information.
-   * This is ignored for distances between objects of different types.
+   * Grouping is only performed when the distances structure contains latencies,
+   * and when all objects are of the same type.
    * \hideinitializer
    */
   HWLOC_DISTANCES_ADD_FLAG_GROUP = (1UL<<0),

src/3rdparty/hwloc/include/hwloc/helper.h

@@ -1,6 +1,6 @@
 /*
  * Copyright © 2009 CNRS
- * Copyright © 2009-2023 Inria.  All rights reserved.
+ * Copyright © 2009-2024 Inria.  All rights reserved.
  * Copyright © 2009-2012 Université Bordeaux
  * Copyright © 2009-2010 Cisco Systems, Inc.  All rights reserved.
  * See COPYING in top-level directory.
@@ -946,6 +946,14 @@ enum hwloc_distrib_flags_e {
  *
  * \return 0 on success, -1 on error.
  *
+ * \note On hybrid CPUs (or asymmetric platforms), distribution may be suboptimal
+ * since the number of cores or PUs inside packages or below caches may vary
+ * (the top-down recursive partitioning ignores these numbers until reaching their levels).
+ * Hence it is recommended to distribute only inside a single homogeneous domain.
+ * For instance on a CPU with energy-efficient E-cores and high-performance P-cores,
+ * one should distribute separately N tasks on E-cores and M tasks on P-cores
+ * instead of trying to distribute directly M+N tasks on the entire CPUs.
+ *
  * \note This function requires the \p roots objects to have a CPU set.
  */
 static __hwloc_inline int
@@ -960,7 +968,7 @@ hwloc_distrib(hwloc_topology_t topology,
   unsigned given, givenweight;
   hwloc_cpuset_t *cpusetp = set;
 
-  if (flags & ~HWLOC_DISTRIB_FLAG_REVERSE) {
+  if (!n || (flags & ~HWLOC_DISTRIB_FLAG_REVERSE)) {
     errno = EINVAL;
     return -1;
   }

src/3rdparty/hwloc/include/hwloc/memattrs.h

@@ -1,5 +1,5 @@
 /*
- * Copyright © 2019-2023 Inria.  All rights reserved.
+ * Copyright © 2019-2024 Inria.  All rights reserved.
  * See COPYING in top-level directory.
  */
 
@@ -69,7 +69,10 @@ extern "C" {
  * @{
  */
 
-/** \brief Memory node attributes. */
+/** \brief Predefined memory attribute IDs.
+ * See ::hwloc_memattr_id_t for the generic definition of IDs
+ * for predefined or custom attributes.
+ */
 enum hwloc_memattr_id_e {
   /** \brief
    * The \"Capacity\" is returned in bytes (local_memory attribute in objects).
@@ -78,6 +81,8 @@ enum hwloc_memattr_id_e {
    *
    * No initiator is involved when looking at this attribute.
    * The corresponding attribute flags are ::HWLOC_MEMATTR_FLAG_HIGHER_FIRST.
+   *
+   * Capacity values may not be modified using hwloc_memattr_set_value().
    * \hideinitializer
    */
   HWLOC_MEMATTR_ID_CAPACITY = 0,
@@ -93,6 +98,8 @@ enum hwloc_memattr_id_e {
    *
    * No initiator is involved when looking at this attribute.
    * The corresponding attribute flags are ::HWLOC_MEMATTR_FLAG_HIGHER_FIRST.
+
+   * Locality values may not be modified using hwloc_memattr_set_value().
    * \hideinitializer
    */
   HWLOC_MEMATTR_ID_LOCALITY = 1,
@@ -173,11 +180,19 @@ enum hwloc_memattr_id_e {
 
   /* TODO persistence? */
 
-  HWLOC_MEMATTR_ID_MAX /**< \private Sentinel value */
+  HWLOC_MEMATTR_ID_MAX /**< \private
+                        * Sentinel value for predefined attributes.
+                        * Dynamically registered custom attributes start here.
+                        */
 };
 
 /** \brief A memory attribute identifier.
- * May be either one of ::hwloc_memattr_id_e or a new id returned by hwloc_memattr_register().
+ *
+ * hwloc predefines some commonly-used attributes in ::hwloc_memattr_id_e.
+ * One may then dynamically register custom ones with hwloc_memattr_register(),
+ * they will be assigned IDs immediately after the predefined ones.
+ * See \ref hwlocality_memattrs_manage for more information about
+ * existing attribute IDs.
  */
 typedef unsigned hwloc_memattr_id_t;
 
@@ -283,6 +298,10 @@ hwloc_get_local_numanode_objs(hwloc_topology_t topology,
  * (it does not have the flag ::HWLOC_MEMATTR_FLAG_NEED_INITIATOR),
  * location \p initiator is ignored and may be \c NULL.
  *
+ * \p target_node cannot be \c NULL. If \p attribute is ::HWLOC_MEMATTR_ID_CAPACITY,
+ * \p target_node must be a NUMA node. If it is ::HWLOC_MEMATTR_ID_LOCALITY,
+ * \p target_node must have a CPU set.
+ *
  * \p flags must be \c 0 for now.
  *
  * \return 0 on success.
@@ -352,6 +371,8 @@ hwloc_memattr_get_best_target(hwloc_topology_t topology,
  * The returned initiator should not be modified or freed,
  * it belongs to the topology.
  *
+ * \p target_node cannot be \c NULL.
+ *
  * \p flags must be \c 0 for now.
  *
  * \return 0 on success.
@@ -362,100 +383,10 @@ hwloc_memattr_get_best_target(hwloc_topology_t topology,
 HWLOC_DECLSPEC int
 hwloc_memattr_get_best_initiator(hwloc_topology_t topology,
                                  hwloc_memattr_id_t attribute,
-                                 hwloc_obj_t target,
+                                 hwloc_obj_t target_node,
                                  unsigned long flags,
                                  struct hwloc_location *best_initiator, hwloc_uint64_t *value);
 
-/** @} */
-
-
-/** \defgroup hwlocality_memattrs_manage Managing memory attributes
- * @{
- */
-
-/** \brief Return the name of a memory attribute.
- *
- * \return 0 on success.
- * \return -1 with errno set to \c EINVAL if the attribute does not exist.
- */
-HWLOC_DECLSPEC int
-hwloc_memattr_get_name(hwloc_topology_t topology,
-                       hwloc_memattr_id_t attribute,
-                       const char **name);
-
-/** \brief Return the flags of the given attribute.
- *
- * Flags are a OR'ed set of ::hwloc_memattr_flag_e.
- *
- * \return 0 on success.
- * \return -1 with errno set to \c EINVAL if the attribute does not exist.
- */
-HWLOC_DECLSPEC int
-hwloc_memattr_get_flags(hwloc_topology_t topology,
-                        hwloc_memattr_id_t attribute,
-                        unsigned long *flags);
-
-/** \brief Memory attribute flags.
- * Given to hwloc_memattr_register() and returned by hwloc_memattr_get_flags().
- */
-enum hwloc_memattr_flag_e {
-  /** \brief The best nodes for this memory attribute are those with the higher values.
-   * For instance Bandwidth.
-   */
-  HWLOC_MEMATTR_FLAG_HIGHER_FIRST = (1UL<<0),
-  /** \brief The best nodes for this memory attribute are those with the lower values.
-   * For instance Latency.
-   */
-  HWLOC_MEMATTR_FLAG_LOWER_FIRST = (1UL<<1),
-  /** \brief The value returned for this memory attribute depends on the given initiator.
-   * For instance Bandwidth and Latency, but not Capacity.
-   */
-  HWLOC_MEMATTR_FLAG_NEED_INITIATOR = (1UL<<2)
-};
-
-/** \brief Register a new memory attribute.
- *
- * Add a specific memory attribute that is not defined in ::hwloc_memattr_id_e.
- * Flags are a OR'ed set of ::hwloc_memattr_flag_e. It must contain at least
- * one of ::HWLOC_MEMATTR_FLAG_HIGHER_FIRST or ::HWLOC_MEMATTR_FLAG_LOWER_FIRST.
- *
- * \return 0 on success.
- * \return -1 with errno set to \c EBUSY if another attribute already uses this name.
- */
-HWLOC_DECLSPEC int
-hwloc_memattr_register(hwloc_topology_t topology,
-                       const char *name,
-                       unsigned long flags,
-                       hwloc_memattr_id_t *id);
-
-/** \brief Set an attribute value for a specific target NUMA node.
- *
- * If the attribute does not relate to a specific initiator
- * (it does not have the flag ::HWLOC_MEMATTR_FLAG_NEED_INITIATOR),
- * location \p initiator is ignored and may be \c NULL.
- *
- * The initiator will be copied into the topology,
- * the caller should free anything allocated to store the initiator,
- * for instance the cpuset.
- *
- * \p flags must be \c 0 for now.
- *
- * \note The initiator \p initiator should be of type ::HWLOC_LOCATION_TYPE_CPUSET
- * when referring to accesses performed by CPU cores.
- * ::HWLOC_LOCATION_TYPE_OBJECT is currently unused internally by hwloc,
- * but users may for instance use it to provide custom information about
- * host memory accesses performed by GPUs.
- *
- * \return 0 on success or -1 on error.
- */
-HWLOC_DECLSPEC int
-hwloc_memattr_set_value(hwloc_topology_t topology,
-                        hwloc_memattr_id_t attribute,
-                        hwloc_obj_t target_node,
-                        struct hwloc_location *initiator,
-                        unsigned long flags,
-                        hwloc_uint64_t value);
-
 /** \brief Return the target NUMA nodes that have some values for a given attribute.
  *
  * Return targets for the given attribute in the \p targets array
@@ -519,6 +450,8 @@ hwloc_memattr_get_targets(hwloc_topology_t topology,
  * The returned initiators should not be modified or freed,
  * they belong to the topology.
  *
+ * \p target_node cannot be \c NULL.
+ *
  * \p flags must be \c 0 for now.
  *
  * If the attribute does not relate to a specific initiator
@@ -538,6 +471,131 @@ hwloc_memattr_get_initiators(hwloc_topology_t topology,
                              hwloc_obj_t target_node,
                              unsigned long flags,
                              unsigned *nr, struct hwloc_location *initiators, hwloc_uint64_t *values);
+
+/** @} */
+
+
+/** \defgroup hwlocality_memattrs_manage Managing memory attributes
+ *
+ * Memory attribues are identified by an ID (::hwloc_memattr_id_t)
+ * and a name. hwloc_memattr_get_name() and hwloc_memattr_get_by_name()
+ * convert between them (or return error if the attribute does not exist).
+ *
+ * The set of valid ::hwloc_memattr_id_t is a contigous set starting at \c 0.
+ * It first contains predefined attributes, as listed
+ * in ::hwloc_memattr_id_e (from \c 0 to \c HWLOC_MEMATTR_ID_MAX-1).
+ * Then custom attributes may be dynamically registered with
+ * hwloc_memattr_register(). They will get the following IDs
+ * (\c HWLOC_MEMATTR_ID_MAX for the first one, etc.).
+ *
+ * To iterate over all valid attributes
+ * (either predefined or dynamically registered custom ones),
+ * one may iterate over IDs starting from \c 0 until hwloc_memattr_get_name()
+ * or hwloc_memattr_get_flags() returns an error.
+ *
+ * The values for an existing attribute or for custom dynamically registered ones
+ * may be set or modified with hwloc_memattr_set_value().
+ *
+ * @{
+ */
+
+/** \brief Return the name of a memory attribute.
+ *
+ * The output pointer \p name cannot be \c NULL.
+ *
+ * \return 0 on success.
+ * \return -1 with errno set to \c EINVAL if the attribute does not exist.
+ */
+HWLOC_DECLSPEC int
+hwloc_memattr_get_name(hwloc_topology_t topology,
+                       hwloc_memattr_id_t attribute,
+                       const char **name);
+
+/** \brief Return the flags of the given attribute.
+ *
+ * Flags are a OR'ed set of ::hwloc_memattr_flag_e.
+ *
+ * The output pointer \p flags cannot be \c NULL.
+ *
+ * \return 0 on success.
+ * \return -1 with errno set to \c EINVAL if the attribute does not exist.
+ */
+HWLOC_DECLSPEC int
+hwloc_memattr_get_flags(hwloc_topology_t topology,
+                        hwloc_memattr_id_t attribute,
+                        unsigned long *flags);
+
+/** \brief Memory attribute flags.
+ * Given to hwloc_memattr_register() and returned by hwloc_memattr_get_flags().
+ */
+enum hwloc_memattr_flag_e {
+  /** \brief The best nodes for this memory attribute are those with the higher values.
+   * For instance Bandwidth.
+   */
+  HWLOC_MEMATTR_FLAG_HIGHER_FIRST = (1UL<<0),
+  /** \brief The best nodes for this memory attribute are those with the lower values.
+   * For instance Latency.
+   */
+  HWLOC_MEMATTR_FLAG_LOWER_FIRST = (1UL<<1),
+  /** \brief The value returned for this memory attribute depends on the given initiator.
+   * For instance Bandwidth and Latency, but not Capacity.
+   */
+  HWLOC_MEMATTR_FLAG_NEED_INITIATOR = (1UL<<2)
+};
+
+/** \brief Register a new memory attribute.
+ *
+ * Add a new custom memory attribute.
+ * Flags are a OR'ed set of ::hwloc_memattr_flag_e. It must contain one of
+ * ::HWLOC_MEMATTR_FLAG_HIGHER_FIRST or ::HWLOC_MEMATTR_FLAG_LOWER_FIRST but not both.
+ *
+ * The new attribute \p id is immediately after the last existing attribute ID
+ * (which is either the ID of the last registered attribute if any,
+ * or the ID of the last predefined attribute in ::hwloc_memattr_id_e).
+ *
+ * \return 0 on success.
+ * \return -1 with errno set to \c EINVAL if an invalid set of flags is given.
+ * \return -1 with errno set to \c EBUSY if another attribute already uses this name.
+ */
+HWLOC_DECLSPEC int
+hwloc_memattr_register(hwloc_topology_t topology,
+                       const char *name,
+                       unsigned long flags,
+                       hwloc_memattr_id_t *id);
+
+/** \brief Set an attribute value for a specific target NUMA node.
+ *
+ * If the attribute does not relate to a specific initiator
+ * (it does not have the flag ::HWLOC_MEMATTR_FLAG_NEED_INITIATOR),
+ * location \p initiator is ignored and may be \c NULL.
+ *
+ * The initiator will be copied into the topology,
+ * the caller should free anything allocated to store the initiator,
+ * for instance the cpuset.
+ *
+ * \p target_node cannot be \c NULL.
+ *
+ * \p attribute cannot be ::HWLOC_MEMATTR_FLAG_ID_CAPACITY or
+ * ::HWLOC_MEMATTR_FLAG_ID_LOCALITY.
+ *
+ * \p flags must be \c 0 for now.
+ *
+ * \note The initiator \p initiator should be of type ::HWLOC_LOCATION_TYPE_CPUSET
+ * when referring to accesses performed by CPU cores.
+ * ::HWLOC_LOCATION_TYPE_OBJECT is currently unused internally by hwloc,
+ * but users may for instance use it to provide custom information about
+ * host memory accesses performed by GPUs.
+ *
+ * \return 0 on success or -1 on error.
+ */
+HWLOC_DECLSPEC int
+hwloc_memattr_set_value(hwloc_topology_t topology,
+                        hwloc_memattr_id_t attribute,
+                        hwloc_obj_t target_node,
+                        struct hwloc_location *initiator,
+                        unsigned long flags,
+                        hwloc_uint64_t value);
+
 /** @} */
 
 #ifdef __cplusplus

src/3rdparty/hwloc/include/hwloc/opencl.h

@@ -41,6 +41,15 @@ extern "C" {
  */
 /* Copyright (c) 2008-2018 The Khronos Group Inc. */
 
+/* needs "cl_khr_pci_bus_info" device extension, but not strictly required for clGetDeviceInfo() */
+typedef struct {
+    cl_uint pci_domain;
+    cl_uint pci_bus;
+    cl_uint pci_device;
+    cl_uint pci_function;
+} hwloc_cl_device_pci_bus_info_khr;
+#define HWLOC_CL_DEVICE_PCI_BUS_INFO_KHR 0x410F
+
 /* needs "cl_amd_device_attribute_query" device extension, but not strictly required for clGetDeviceInfo() */
 #define HWLOC_CL_DEVICE_TOPOLOGY_AMD 0x4037
 typedef union {
@@ -78,9 +87,19 @@ hwloc_opencl_get_device_pci_busid(cl_device_id device,
                                unsigned *domain, unsigned *bus, unsigned *dev, unsigned *func)
 {
 	hwloc_cl_device_topology_amd amdtopo;
+	hwloc_cl_device_pci_bus_info_khr khrbusinfo;
 	cl_uint nvbus, nvslot, nvdomain;
 	cl_int clret;
 
+	clret = clGetDeviceInfo(device, HWLOC_CL_DEVICE_PCI_BUS_INFO_KHR, sizeof(khrbusinfo), &khrbusinfo, NULL);
+	if (CL_SUCCESS == clret) {
+		*domain = (unsigned) khrbusinfo.pci_domain;
+		*bus = (unsigned) khrbusinfo.pci_bus;
+		*dev = (unsigned) khrbusinfo.pci_device;
+		*func = (unsigned) khrbusinfo.pci_function;
+		return 0;
+	}
+
 	clret = clGetDeviceInfo(device, HWLOC_CL_DEVICE_TOPOLOGY_AMD, sizeof(amdtopo), &amdtopo, NULL);
 	if (CL_SUCCESS == clret
 	    && HWLOC_CL_DEVICE_TOPOLOGY_TYPE_PCIE_AMD == amdtopo.raw.type) {

src/3rdparty/hwloc/include/hwloc/plugins.h

@@ -1,5 +1,5 @@
 /*
- * Copyright © 2013-2022 Inria.  All rights reserved.
+ * Copyright © 2013-2024 Inria.  All rights reserved.
  * Copyright © 2016 Cisco Systems, Inc.  All rights reserved.
  * See COPYING in top-level directory.
  */
@@ -645,6 +645,19 @@ HWLOC_DECLSPEC struct hwloc_obj * hwloc_pci_find_parent_by_busid(struct hwloc_to
  */
 HWLOC_DECLSPEC struct hwloc_obj * hwloc_pci_find_by_busid(struct hwloc_topology *topology, unsigned domain, unsigned bus, unsigned dev, unsigned func);
 
+
+/** @} */
+
+
+
+
+/** \defgroup hwlocality_components_distances Components and Plugins: distances
+ *
+ * \note These structures and functions may change when ::HWLOC_COMPONENT_ABI is modified.
+ *
+ * @{
+ */
+
 /** \brief Handle to a new distances structure during its addition to the topology. */
 typedef void * hwloc_backend_distances_add_handle_t;
 

src/3rdparty/hwloc/include/hwloc/rename.h

@@ -1,6 +1,6 @@
 /*
  * Copyright © 2009-2011 Cisco Systems, Inc.  All rights reserved.
- * Copyright © 2010-2022 Inria.  All rights reserved.
+ * Copyright © 2010-2024 Inria.  All rights reserved.
  * See COPYING in top-level directory.
  */
 
@@ -210,6 +210,7 @@ extern "C" {
 
 #define hwloc_obj_get_info_by_name HWLOC_NAME(obj_get_info_by_name)
 #define hwloc_obj_add_info HWLOC_NAME(obj_add_info)
+#define hwloc_obj_set_subtype HWLOC_NAME(obj_set_subtype)
 
 #define HWLOC_CPUBIND_PROCESS HWLOC_NAME_CAPS(CPUBIND_PROCESS)
 #define HWLOC_CPUBIND_THREAD HWLOC_NAME_CAPS(CPUBIND_THREAD)
@@ -232,6 +233,7 @@ extern "C" {
 #define HWLOC_MEMBIND_FIRSTTOUCH HWLOC_NAME_CAPS(MEMBIND_FIRSTTOUCH)
 #define HWLOC_MEMBIND_BIND HWLOC_NAME_CAPS(MEMBIND_BIND)
 #define HWLOC_MEMBIND_INTERLEAVE HWLOC_NAME_CAPS(MEMBIND_INTERLEAVE)
+#define HWLOC_MEMBIND_WEIGHTED_INTERLEAVE HWLOC_NAME_CAPS(MEMBIND_WEIGHTED_INTERLEAVE)
 #define HWLOC_MEMBIND_NEXTTOUCH HWLOC_NAME_CAPS(MEMBIND_NEXTTOUCH)
 #define HWLOC_MEMBIND_MIXED HWLOC_NAME_CAPS(MEMBIND_MIXED)
 
@@ -560,6 +562,7 @@ extern "C" {
 
 /* opencl.h */
 
+#define hwloc_cl_device_pci_bus_info_khr HWLOC_NAME(cl_device_pci_bus_info_khr)
 #define hwloc_cl_device_topology_amd HWLOC_NAME(cl_device_topology_amd)
 #define hwloc_opencl_get_device_pci_busid HWLOC_NAME(opencl_get_device_pci_ids)
 #define hwloc_opencl_get_device_cpuset HWLOC_NAME(opencl_get_device_cpuset)
@@ -715,6 +718,8 @@ extern "C" {
 #define hwloc__obj_type_is_dcache HWLOC_NAME(_obj_type_is_dcache)
 #define hwloc__obj_type_is_icache HWLOC_NAME(_obj_type_is_icache)
 
+#define hwloc__pci_link_speed HWLOC_NAME(_pci_link_speed)
+
 /* private/cpuid-x86.h */
 
 #define hwloc_have_x86_cpuid HWLOC_NAME(have_x86_cpuid)

src/3rdparty/hwloc/include/private/autogen/config.h

@@ -1,6 +1,6 @@
 /*
  * Copyright © 2009, 2011, 2012 CNRS.  All rights reserved.
- * Copyright © 2009-2021 Inria.  All rights reserved.
+ * Copyright © 2009-2020 Inria.  All rights reserved.
  * Copyright © 2009, 2011, 2012, 2015 Université Bordeaux.  All rights reserved.
  * Copyright © 2009-2020 Cisco Systems, Inc.  All rights reserved.
  * $COPYRIGHT$
@@ -17,6 +17,10 @@
 
 #define HWLOC_HAVE_MSVC_CPUIDEX 1
 
+/* #undef HAVE_MKSTEMP */
+
+#define HWLOC_HAVE_X86_CPUID 1
+
 /* Define to 1 if the system has the type `CACHE_DESCRIPTOR'. */
 #define HAVE_CACHE_DESCRIPTOR 0
 
@@ -128,8 +132,7 @@
 #define HAVE_DECL__SC_PAGE_SIZE 0
 
 /* Define to 1 if you have the <dirent.h> header file. */
-/* #define HAVE_DIRENT_H 1 */
-#undef HAVE_DIRENT_H
+/* #undef HAVE_DIRENT_H */
 
 /* Define to 1 if you have the <dlfcn.h> header file. */
 /* #undef HAVE_DLFCN_H */
@@ -282,7 +285,7 @@
 #define HAVE_STRING_H 1
 
 /* Define to 1 if you have the `strncasecmp' function. */
-#define HAVE_STRNCASECMP 1
+/* #undef HAVE_STRNCASECMP */
 
 /* Define to '1' if sysctl is present and usable */
 /* #undef HAVE_SYSCTL */
@@ -323,8 +326,7 @@
 /* #undef HAVE_UNAME */
 
 /* Define to 1 if you have the <unistd.h> header file. */
-/* #define HAVE_UNISTD_H 1 */
-#undef HAVE_UNISTD_H
+/* #undef HAVE_UNISTD_H */
 
 /* Define to 1 if you have the `uselocale' function. */
 /* #undef HAVE_USELOCALE */
@@ -659,7 +661,7 @@
 #define hwloc_pid_t HANDLE
 
 /* Define this to either strncasecmp or strncmp */
-#define hwloc_strncasecmp strncasecmp
+/* #undef hwloc_strncasecmp */
 
 /* Define this to the thread ID type */
 #define hwloc_thread_t HANDLE

src/3rdparty/hwloc/include/private/cpuid-x86.h

@@ -11,6 +11,22 @@
 #ifndef HWLOC_PRIVATE_CPUID_X86_H
 #define HWLOC_PRIVATE_CPUID_X86_H
 
+/* A macro for annotating memory as uninitialized when building with MSAN
+ * (and otherwise having no effect). See below for why this is used with
+ * our custom assembly.
+ */
+#ifdef __has_feature
+#define HWLOC_HAS_FEATURE(name) __has_feature(name)
+#else
+#define HWLOC_HAS_FEATURE(name) 0
+#endif
+#if HWLOC_HAS_FEATURE(memory_sanitizer) || defined(MEMORY_SANITIZER)
+#include <sanitizer/msan_interface.h>
+#define HWLOC_ANNOTATE_MEMORY_IS_INITIALIZED(ptr, len) __msan_unpoison(ptr, len)
+#else
+#define HWLOC_ANNOTATE_MEMORY_IS_INITIALIZED(ptr, len)
+#endif
+
 #if (defined HWLOC_X86_32_ARCH) && (!defined HWLOC_HAVE_MSVC_CPUIDEX)
 static __hwloc_inline int hwloc_have_x86_cpuid(void)
 {
@@ -71,12 +87,18 @@ static __hwloc_inline void hwloc_x86_cpuid(unsigned *eax, unsigned *ebx, unsigne
   "movl %k2,%1\n\t"
   : "+a" (*eax), "=m" (*ebx), "=&r"(sav_rbx),
     "+c" (*ecx), "=&d" (*edx));
+  /* MSAN does not recognize the effect of the above assembly on the memory operand
+   * (`"=m"(*ebx)`). This may get improved in MSAN at some point in the future, e.g.
+   * see https://github.com/llvm/llvm-project/pull/77393. */
+  HWLOC_ANNOTATE_MEMORY_IS_INITIALIZED(ebx, sizeof *ebx);
 #elif defined(HWLOC_X86_32_ARCH)
   __asm__(
   "mov %%ebx,%1\n\t"
   "cpuid\n\t"
   "xchg %%ebx,%1\n\t"
   : "+a" (*eax), "=&SD" (*ebx), "+c" (*ecx), "=&d" (*edx));
+  /* See above. */
+  HWLOC_ANNOTATE_MEMORY_IS_INITIALIZED(ebx, sizeof *ebx);
 #else
 #error unknown architecture
 #endif

src/3rdparty/hwloc/include/private/misc.h

@@ -1,6 +1,6 @@
 /*
  * Copyright © 2009 CNRS
- * Copyright © 2009-2019 Inria.  All rights reserved.
+ * Copyright © 2009-2024 Inria.  All rights reserved.
  * Copyright © 2009-2012 Université Bordeaux
  * Copyright © 2011 Cisco Systems, Inc.  All rights reserved.
  * See COPYING in top-level directory.
@@ -573,4 +573,35 @@ typedef SSIZE_T ssize_t;
 #  endif
 #endif
 
+static __inline float
+hwloc__pci_link_speed(unsigned generation, unsigned lanes)
+{
+  float lanespeed;
+  /*
+   * These are single-direction bandwidths only.
+   *
+   * Gen1 used NRZ with 8/10 encoding.
+   * PCIe Gen1 = 2.5GT/s signal-rate per lane x 8/10    =  0.25GB/s data-rate per lane
+   * PCIe Gen2 = 5  GT/s signal-rate per lane x 8/10    =  0.5 GB/s data-rate per lane
+   * Gen3 switched to NRZ with 128/130 encoding.
+   * PCIe Gen3 = 8  GT/s signal-rate per lane x 128/130 =  1   GB/s data-rate per lane
+   * PCIe Gen4 = 16 GT/s signal-rate per lane x 128/130 =  2   GB/s data-rate per lane
+   * PCIe Gen5 = 32 GT/s signal-rate per lane x 128/130 =  4   GB/s data-rate per lane
+   * Gen6 switched to PAM with with 242/256 FLIT (242B payload protected by 8B CRC + 6B FEC).
+   * PCIe Gen6 = 64 GT/s signal-rate per lane x 242/256 =  8   GB/s data-rate per lane
+   * PCIe Gen7 = 128GT/s signal-rate per lane x 242/256 = 16   GB/s data-rate per lane
+   */
+
+  /* lanespeed in Gbit/s */
+  if (generation <= 2)
+    lanespeed = 2.5f * generation * 0.8f;
+  else if (generation <= 5)
+    lanespeed = 8.0f * (1<<(generation-3)) * 128/130;
+  else
+    lanespeed = 8.0f * (1<<(generation-3)) * 242/256; /* assume Gen8 will be 256 GT/s and so on */
+
+  /* linkspeed in GB/s */
+  return lanespeed * lanes / 8;
+}
+
 #endif /* HWLOC_PRIVATE_MISC_H */

src/3rdparty/hwloc/src/bind.c

@@ -1,6 +1,6 @@
 /*
  * Copyright © 2009 CNRS
- * Copyright © 2009-2020 Inria.  All rights reserved.
+ * Copyright © 2009-2024 Inria.  All rights reserved.
  * Copyright © 2009-2010, 2012 Université Bordeaux
  * Copyright © 2011-2015 Cisco Systems, Inc.  All rights reserved.
  * See COPYING in top-level directory.
@@ -287,6 +287,7 @@ static __hwloc_inline int hwloc__check_membind_policy(hwloc_membind_policy_t pol
       || policy == HWLOC_MEMBIND_FIRSTTOUCH
       || policy == HWLOC_MEMBIND_BIND
       || policy == HWLOC_MEMBIND_INTERLEAVE
+      || policy == HWLOC_MEMBIND_WEIGHTED_INTERLEAVE
       || policy == HWLOC_MEMBIND_NEXTTOUCH)
     return 0;
   return -1;

src/3rdparty/hwloc/src/bitmap.c

@@ -1,6 +1,6 @@
 /*
  * Copyright © 2009 CNRS
- * Copyright © 2009-2020 Inria.  All rights reserved.
+ * Copyright © 2009-2024 Inria.  All rights reserved.
  * Copyright © 2009-2011 Université Bordeaux
  * Copyright © 2009-2011 Cisco Systems, Inc.  All rights reserved.
  * See COPYING in top-level directory.
@@ -245,6 +245,7 @@ int hwloc_bitmap_copy(struct hwloc_bitmap_s * dst, const struct hwloc_bitmap_s *
 /* Strings always use 32bit groups */
 #define HWLOC_PRIxSUBBITMAP		"%08lx"
 #define HWLOC_BITMAP_SUBSTRING_SIZE	32
+#define HWLOC_BITMAP_SUBSTRING_FULL_VALUE 0xFFFFFFFFUL
 #define HWLOC_BITMAP_SUBSTRING_LENGTH	(HWLOC_BITMAP_SUBSTRING_SIZE/4)
 #define HWLOC_BITMAP_STRING_PER_LONG	(HWLOC_BITS_PER_LONG/HWLOC_BITMAP_SUBSTRING_SIZE)
 
@@ -261,6 +262,7 @@ int hwloc_bitmap_snprintf(char * __hwloc_restrict buf, size_t buflen, const stru
   const unsigned long accum_mask = ~0UL;
 #else /* HWLOC_BITS_PER_LONG != HWLOC_BITMAP_SUBSTRING_SIZE */
   const unsigned long accum_mask = ((1UL << HWLOC_BITMAP_SUBSTRING_SIZE) - 1) << (HWLOC_BITS_PER_LONG - HWLOC_BITMAP_SUBSTRING_SIZE);
+  int merge_with_infinite_prefix = 0;
 #endif /* HWLOC_BITS_PER_LONG != HWLOC_BITMAP_SUBSTRING_SIZE */
 
   HWLOC__BITMAP_CHECK(set);
@@ -279,6 +281,9 @@ int hwloc_bitmap_snprintf(char * __hwloc_restrict buf, size_t buflen, const stru
       res = size>0 ? (int)size - 1 : 0;
     tmp += res;
     size -= res;
+#if HWLOC_BITS_PER_LONG > HWLOC_BITMAP_SUBSTRING_SIZE
+    merge_with_infinite_prefix = 1;
+#endif
   }
 
   i=(int) set->ulongs_count-1;
@@ -294,16 +299,24 @@ int hwloc_bitmap_snprintf(char * __hwloc_restrict buf, size_t buflen, const stru
   }
 
   while (i>=0 || accumed) {
+    unsigned long value;
+
     /* Refill accumulator */
     if (!accumed) {
       accum = set->ulongs[i--];
       accumed = HWLOC_BITS_PER_LONG;
     }
+    value = (accum & accum_mask) >> (HWLOC_BITS_PER_LONG - HWLOC_BITMAP_SUBSTRING_SIZE);
 
-    if (accum & accum_mask) {
+#if HWLOC_BITS_PER_LONG > HWLOC_BITMAP_SUBSTRING_SIZE
+    if (merge_with_infinite_prefix && value == HWLOC_BITMAP_SUBSTRING_FULL_VALUE) {
+      /* first full subbitmap merged with infinite prefix */
+      res = 0;
+    } else
+#endif
+    if (value) {
       /* print the whole subset if not empty */
-        res = hwloc_snprintf(tmp, size, needcomma ? ",0x" HWLOC_PRIxSUBBITMAP : "0x" HWLOC_PRIxSUBBITMAP,
-		     (accum & accum_mask) >> (HWLOC_BITS_PER_LONG - HWLOC_BITMAP_SUBSTRING_SIZE));
+      res = hwloc_snprintf(tmp, size, needcomma ? ",0x" HWLOC_PRIxSUBBITMAP : "0x" HWLOC_PRIxSUBBITMAP, value);
       needcomma = 1;
     } else if (i == -1 && accumed == HWLOC_BITMAP_SUBSTRING_SIZE) {
       /* print a single 0 to mark the last subset */
@@ -323,6 +336,7 @@ int hwloc_bitmap_snprintf(char * __hwloc_restrict buf, size_t buflen, const stru
 #else
     accum <<= HWLOC_BITMAP_SUBSTRING_SIZE;
     accumed -= HWLOC_BITMAP_SUBSTRING_SIZE;
+    merge_with_infinite_prefix = 0;
 #endif
 
     if (res >= size)
@@ -363,6 +377,7 @@ int hwloc_bitmap_sscanf(struct hwloc_bitmap_s *set, const char * __hwloc_restric
   const char * current = string;
   unsigned long accum = 0;
   int count = 0;
+  int ulongcount;
   int infinite = 0;
 
   /* count how many substrings there are */
@@ -383,9 +398,20 @@ int hwloc_bitmap_sscanf(struct hwloc_bitmap_s *set, const char * __hwloc_restric
     count--;
   }
 
-  if (hwloc_bitmap_reset_by_ulongs(set, (count + HWLOC_BITMAP_STRING_PER_LONG - 1) / HWLOC_BITMAP_STRING_PER_LONG) < 0)
+  ulongcount = (count + HWLOC_BITMAP_STRING_PER_LONG - 1) / HWLOC_BITMAP_STRING_PER_LONG;
+  if (hwloc_bitmap_reset_by_ulongs(set, ulongcount) < 0)
     return -1;
-  set->infinite = 0;
+
+  set->infinite = 0; /* will be updated later */
+
+#if HWLOC_BITS_PER_LONG != HWLOC_BITMAP_SUBSTRING_SIZE
+  if (infinite && (count % HWLOC_BITMAP_STRING_PER_LONG) != 0) {
+    /* accumulate substrings of the first ulong that are hidden in the infinite prefix */
+    int i;
+    for(i = (count % HWLOC_BITMAP_STRING_PER_LONG); i < HWLOC_BITMAP_STRING_PER_LONG; i++)
+      accum |= (HWLOC_BITMAP_SUBSTRING_FULL_VALUE << (i*HWLOC_BITMAP_SUBSTRING_SIZE));
+  }
+#endif
 
   while (*current != '\0') {
     unsigned long val;
@@ -544,6 +570,9 @@ int hwloc_bitmap_taskset_snprintf(char * __hwloc_restrict buf, size_t buflen, co
   ssize_t size = buflen;
   char *tmp = buf;
   int res, ret = 0;
+#if HWLOC_BITS_PER_LONG == 64
+  int merge_with_infinite_prefix = 0;
+#endif
   int started = 0;
   int i;
 
@@ -563,6 +592,9 @@ int hwloc_bitmap_taskset_snprintf(char * __hwloc_restrict buf, size_t buflen, co
       res = size>0 ? (int)size - 1 : 0;
     tmp += res;
     size -= res;
+#if HWLOC_BITS_PER_LONG == 64
+    merge_with_infinite_prefix = 1;
+#endif
   }
 
   i=set->ulongs_count-1;
@@ -582,7 +614,11 @@ int hwloc_bitmap_taskset_snprintf(char * __hwloc_restrict buf, size_t buflen, co
     if (started) {
       /* print the whole subset */
 #if HWLOC_BITS_PER_LONG == 64
+      if (merge_with_infinite_prefix && (val & 0xffffffff00000000UL) == 0xffffffff00000000UL) {
+        res = hwloc_snprintf(tmp, size, "%08lx", val & 0xffffffffUL);
+      } else  {
         res = hwloc_snprintf(tmp, size, "%016lx", val);
+      }
 #else
       res = hwloc_snprintf(tmp, size, "%08lx", val);
 #endif
@@ -599,6 +635,9 @@ int hwloc_bitmap_taskset_snprintf(char * __hwloc_restrict buf, size_t buflen, co
       res = size>0 ? (int)size - 1 : 0;
     tmp += res;
     size -= res;
+#if HWLOC_BITS_PER_LONG == 64
+    merge_with_infinite_prefix = 0;
+#endif
   }
 
   /* if didn't display anything, display 0x0 */
@@ -679,6 +718,10 @@ int hwloc_bitmap_taskset_sscanf(struct hwloc_bitmap_s *set, const char * __hwloc
       goto failed;
 
     set->ulongs[count-1] = val;
+    if (infinite && tmpchars != HWLOC_BITS_PER_LONG/4) {
+      /* infinite prefix with partial substring, fill remaining bits */
+      set->ulongs[count-1] |= (~0ULL)<<(4*tmpchars);
+    }
 
     current += tmpchars;
     chars -= tmpchars;

src/3rdparty/hwloc/src/cpukinds.c

@@ -1,5 +1,5 @@
 /*
- * Copyright © 2020-2022 Inria.  All rights reserved.
+ * Copyright © 2020-2024 Inria.  All rights reserved.
  * See COPYING in top-level directory.
  */
 
@@ -50,6 +50,7 @@ hwloc_internal_cpukinds_dup(hwloc_topology_t new, hwloc_topology_t old)
     return -1;
   new->cpukinds = kinds;
   new->nr_cpukinds = old->nr_cpukinds;
+  new->nr_cpukinds_allocated = old->nr_cpukinds;
   memcpy(kinds, old->cpukinds, old->nr_cpukinds * sizeof(*kinds));
 
   for(i=0;i<old->nr_cpukinds; i++) {

src/3rdparty/hwloc/src/distances.c

@@ -1,5 +1,5 @@
 /*
- * Copyright © 2010-2022 Inria.  All rights reserved.
+ * Copyright © 2010-2024 Inria.  All rights reserved.
  * Copyright © 2011-2012 Université Bordeaux
  * Copyright © 2011 Cisco Systems, Inc.  All rights reserved.
  * See COPYING in top-level directory.
@@ -624,8 +624,8 @@ void * hwloc_distances_add_create(hwloc_topology_t topology,
     return NULL;
   }
   if ((kind & ~HWLOC_DISTANCES_KIND_ALL)
-      || hwloc_weight_long(kind & HWLOC_DISTANCES_KIND_FROM_ALL) != 1
-      || hwloc_weight_long(kind & HWLOC_DISTANCES_KIND_MEANS_ALL) != 1) {
+      || hwloc_weight_long(kind & HWLOC_DISTANCES_KIND_FROM_ALL) > 1
+      || hwloc_weight_long(kind & HWLOC_DISTANCES_KIND_MEANS_ALL) > 1) {
     errno = EINVAL;
     return NULL;
   }

src/3rdparty/hwloc/src/memattrs.c

@@ -1,5 +1,5 @@
 /*
- * Copyright © 2020-2023 Inria.  All rights reserved.
+ * Copyright © 2020-2024 Inria.  All rights reserved.
  * See COPYING in top-level directory.
  */
 
@@ -14,13 +14,26 @@
  */
 
 static __hwloc_inline
-hwloc_uint64_t hwloc__memattr_get_convenience_value(hwloc_memattr_id_t id,
-                                                    hwloc_obj_t node)
+int hwloc__memattr_get_convenience_value(hwloc_memattr_id_t id,
+                                         hwloc_obj_t node,
+                                         hwloc_uint64_t *valuep)
 {
-  if (id == HWLOC_MEMATTR_ID_CAPACITY)
-    return node->attr->numanode.local_memory;
-  else if (id == HWLOC_MEMATTR_ID_LOCALITY)
-    return hwloc_bitmap_weight(node->cpuset);
+  if (id == HWLOC_MEMATTR_ID_CAPACITY) {
+    if (node->type != HWLOC_OBJ_NUMANODE) {
+      errno = EINVAL;
+      return -1;
+    }
+    *valuep = node->attr->numanode.local_memory;
+    return 0;
+  }
+  else if (id == HWLOC_MEMATTR_ID_LOCALITY) {
+    if (!node->cpuset) {
+      errno = EINVAL;
+      return -1;
+    }
+    *valuep = hwloc_bitmap_weight(node->cpuset);
+    return 0;
+  }
   else
     assert(0);
   return 0; /* shut up the compiler */
@@ -622,7 +635,7 @@ hwloc_memattr_get_targets(hwloc_topology_t topology,
       if (found<max) {
         targets[found] = node;
         if (values)
-          values[found] = hwloc__memattr_get_convenience_value(id, node);
+          hwloc__memattr_get_convenience_value(id, node, &values[found]);
       }
       found++;
     }
@@ -748,7 +761,7 @@ hwloc_memattr_get_initiators(hwloc_topology_t topology,
   struct hwloc_internal_memattr_target_s *imtg;
   unsigned i, max;
 
-  if (flags) {
+  if (flags || !target_node) {
     errno = EINVAL;
     return -1;
   }
@@ -810,7 +823,7 @@ hwloc_memattr_get_value(hwloc_topology_t topology,
   struct hwloc_internal_memattr_s *imattr;
   struct hwloc_internal_memattr_target_s *imtg;
 
-  if (flags) {
+  if (flags || !target_node) {
     errno = EINVAL;
     return -1;
   }
@@ -823,8 +836,7 @@ hwloc_memattr_get_value(hwloc_topology_t topology,
 
   if (imattr->iflags & HWLOC_IMATTR_FLAG_CONVENIENCE) {
     /* convenience attributes */
-    *valuep = hwloc__memattr_get_convenience_value(id, target_node);
-    return 0;
+    return hwloc__memattr_get_convenience_value(id, target_node, valuep);
   }
 
   /* normal attributes */
@@ -936,7 +948,7 @@ hwloc_memattr_set_value(hwloc_topology_t topology,
 {
   struct hwloc_internal_location_s iloc, *ilocp;
 
-  if (flags) {
+  if (flags || !target_node) {
     errno = EINVAL;
     return -1;
   }
@@ -1007,10 +1019,10 @@ hwloc_memattr_get_best_target(hwloc_topology_t topology,
     /* convenience attributes */
     for(j=0; ; j++) {
       hwloc_obj_t node = hwloc_get_obj_by_type(topology, HWLOC_OBJ_NUMANODE, j);
-      hwloc_uint64_t value;
+      hwloc_uint64_t value = 0;
       if (!node)
         break;
-      value = hwloc__memattr_get_convenience_value(id, node);
+      hwloc__memattr_get_convenience_value(id, node, &value);
       hwloc__update_best_target(&best, &best_value, &found,
                                 node, value,
                                 imattr->flags & HWLOC_MEMATTR_FLAG_HIGHER_FIRST);
@@ -1093,7 +1105,7 @@ hwloc_memattr_get_best_initiator(hwloc_topology_t topology,
   int found;
   unsigned i;
 
-  if (flags) {
+  if (flags || !target_node) {
     errno = EINVAL;
     return -1;
   }
@@ -1806,6 +1818,12 @@ hwloc__apply_memory_tiers_subtypes(hwloc_topology_t topology,
       }
     }
   }
+  if (nr_tiers > 1) {
+    hwloc_obj_t root = hwloc_get_root_obj(topology);
+    char tmp[20];
+    snprintf(tmp, sizeof(tmp), "%u", nr_tiers);
+    hwloc__add_info_nodup(&root->infos, &root->infos_count, "MemoryTiersNr", tmp, 1);
+  }
 }
 
 int

src/3rdparty/hwloc/src/pci-common.c

@@ -1,5 +1,5 @@
 /*
- * Copyright © 2009-2022 Inria.  All rights reserved.
+ * Copyright © 2009-2024 Inria.  All rights reserved.
  * See COPYING in top-level directory.
  */
 
@@ -886,36 +886,12 @@ hwloc_pcidisc_find_linkspeed(const unsigned char *config,
 			     unsigned offset, float *linkspeed)
 {
   unsigned linksta, speed, width;
-  float lanespeed;
 
   memcpy(&linksta, &config[offset + HWLOC_PCI_EXP_LNKSTA], 4);
   speed = linksta & HWLOC_PCI_EXP_LNKSTA_SPEED; /* PCIe generation */
   width = (linksta & HWLOC_PCI_EXP_LNKSTA_WIDTH) >> 4; /* how many lanes */
-  /*
-   * These are single-direction bandwidths only.
-   *
-   * Gen1 used NRZ with 8/10 encoding.
-   * PCIe Gen1 = 2.5GT/s signal-rate per lane x 8/10    =  0.25GB/s data-rate per lane
-   * PCIe Gen2 = 5  GT/s signal-rate per lane x 8/10    =  0.5 GB/s data-rate per lane
-   * Gen3 switched to NRZ with 128/130 encoding.
-   * PCIe Gen3 = 8  GT/s signal-rate per lane x 128/130 =  1   GB/s data-rate per lane
-   * PCIe Gen4 = 16 GT/s signal-rate per lane x 128/130 =  2   GB/s data-rate per lane
-   * PCIe Gen5 = 32 GT/s signal-rate per lane x 128/130 =  4   GB/s data-rate per lane
-   * Gen6 switched to PAM with with 242/256 FLIT (242B payload protected by 8B CRC + 6B FEC).
-   * PCIe Gen6 = 64 GT/s signal-rate per lane x 242/256 =  8   GB/s data-rate per lane
-   * PCIe Gen7 = 128GT/s signal-rate per lane x 242/256 = 16   GB/s data-rate per lane
-   */
 
-  /* lanespeed in Gbit/s */
-  if (speed <= 2)
-    lanespeed = 2.5f * speed * 0.8f;
-  else if (speed <= 5)
-    lanespeed = 8.0f * (1<<(speed-3)) * 128/130;
-  else
-    lanespeed = 8.0f * (1<<(speed-3)) * 242/256; /* assume Gen8 will be 256 GT/s and so on */
-
-  /* linkspeed in GB/s */
-  *linkspeed = lanespeed * width / 8;
+  *linkspeed = hwloc__pci_link_speed(speed, width);
   return 0;
 }
 

src/3rdparty/hwloc/src/topology-windows.c

@@ -1,6 +1,6 @@
 /*
  * Copyright © 2009 CNRS
- * Copyright © 2009-2023 Inria.  All rights reserved.
+ * Copyright © 2009-2024 Inria.  All rights reserved.
  * Copyright © 2009-2012, 2020 Université Bordeaux
  * Copyright © 2011 Cisco Systems, Inc.  All rights reserved.
  * See COPYING in top-level directory.
@@ -220,7 +220,7 @@ static void hwloc_win_get_function_ptrs(void)
 #pragma GCC diagnostic ignored "-Wcast-function-type"
 #endif
 
-    kernel32 = LoadLibrary("kernel32.dll");
+    kernel32 = LoadLibrary(TEXT("kernel32.dll"));
     if (kernel32) {
       GetActiveProcessorGroupCountProc =
 	(PFN_GETACTIVEPROCESSORGROUPCOUNT) GetProcAddress(kernel32, "GetActiveProcessorGroupCount");
@@ -249,12 +249,12 @@ static void hwloc_win_get_function_ptrs(void)
     }
 
     if (!QueryWorkingSetExProc) {
-      HMODULE psapi = LoadLibrary("psapi.dll");
+      HMODULE psapi = LoadLibrary(TEXT("psapi.dll"));
       if (psapi)
         QueryWorkingSetExProc = (PFN_QUERYWORKINGSETEX) GetProcAddress(psapi, "QueryWorkingSetEx");
     }
 
-    ntdll = GetModuleHandle("ntdll");
+    ntdll = GetModuleHandle(TEXT("ntdll"));
     RtlGetVersionProc = (PFN_RTLGETVERSION) GetProcAddress(ntdll, "RtlGetVersion");
 
 #if HWLOC_HAVE_GCC_W_CAST_FUNCTION_TYPE

src/3rdparty/hwloc/src/topology-x86.c

@@ -1,11 +1,11 @@
 /*
- * Copyright © 2010-2023 Inria.  All rights reserved.
+ * Copyright © 2010-2024 Inria.  All rights reserved.
  * Copyright © 2010-2013 Université Bordeaux
  * Copyright © 2010-2011 Cisco Systems, Inc.  All rights reserved.
  * See COPYING in top-level directory.
  *
  *
- * This backend is only used when the operating system does not export
+ * This backend is mostly used when the operating system does not export
  * the necessary hardware topology information to user-space applications.
  * Currently, FreeBSD and NetBSD only add PUs and then fallback to this
  * backend for CPU/Cache discovery.
@@ -15,6 +15,7 @@
  * on various architectures, without having to use this x86-specific code.
  * But this backend is still used after them to annotate some objects with
  * additional details (CPU info in Package, Inclusiveness in Caches).
+ * It may also be enabled manually to work-around bugs in native OS discovery.
  */
 
 #include "private/autogen/config.h"
@@ -487,7 +488,7 @@ static void read_amd_cores_legacy(struct procinfo *infos, struct cpuiddump *src_
 }
 
 /* AMD unit/node from CPUID 0x8000001e leaf (topoext) */
-static void read_amd_cores_topoext(struct hwloc_x86_backend_data_s *data, struct procinfo *infos, unsigned long flags, struct cpuiddump *src_cpuiddump)
+static void read_amd_cores_topoext(struct hwloc_x86_backend_data_s *data, struct procinfo *infos, unsigned long flags __hwloc_attribute_unused, struct cpuiddump *src_cpuiddump)
 {
   unsigned apic_id, nodes_per_proc = 0;
   unsigned eax, ebx, ecx, edx;
@@ -496,7 +497,6 @@ static void read_amd_cores_topoext(struct hwloc_x86_backend_data_s *data, struct
   cpuid_or_from_dump(&eax, &ebx, &ecx, &edx, src_cpuiddump);
   infos->apicid = apic_id = eax;
 
-  if (flags & HWLOC_X86_DISC_FLAG_TOPOEXT_NUMANODES) {
     if (infos->cpufamilynumber == 0x16) {
       /* ecx is reserved */
       infos->ids[NODE] = 0;
@@ -511,7 +511,6 @@ static void read_amd_cores_topoext(struct hwloc_x86_backend_data_s *data, struct
         || (infos->cpufamilynumber == 0x19 && nodes_per_proc > 1)) {
       hwloc_debug("warning: undefined nodes_per_proc value %u, assuming it means %u\n", nodes_per_proc, nodes_per_proc);
     }
-  }
 
   if (infos->cpufamilynumber <= 0x16) { /* topoext appeared in 0x15 and compute-units were only used in 0x15 and 0x16 */
     unsigned cores_per_unit;
@@ -533,9 +532,9 @@ static void read_amd_cores_topoext(struct hwloc_x86_backend_data_s *data, struct
 }
 
 /* Intel core/thread or even die/module/tile from CPUID 0x0b or 0x1f leaves (v1 and v2 extended topology enumeration)
- * or AMD complex/ccd from CPUID 0x80000026 (extended CPU topology)
+ * or AMD core/thread or even complex/ccd from CPUID 0x0b or 0x80000026 (extended CPU topology)
  */
-static void read_extended_topo(struct hwloc_x86_backend_data_s *data, struct procinfo *infos, unsigned leaf, enum cpuid_type cpuid_type, struct cpuiddump *src_cpuiddump)
+static void read_extended_topo(struct hwloc_x86_backend_data_s *data, struct procinfo *infos, unsigned leaf, enum cpuid_type cpuid_type __hwloc_attribute_unused, struct cpuiddump *src_cpuiddump)
 {
   unsigned level, apic_nextshift, apic_type, apic_id = 0, apic_shift = 0, id;
   unsigned threadid __hwloc_attribute_unused = 0; /* shut-up compiler */
@@ -547,20 +546,15 @@ static void read_extended_topo(struct hwloc_x86_backend_data_s *data, struct pro
     eax = leaf;
     cpuid_or_from_dump(&eax, &ebx, &ecx, &edx, src_cpuiddump);
     /* Intel specifies that the 0x0b/0x1f loop should stop when we get "invalid domain" (0 in ecx[8:15])
-     * (if so, we also get 0 in eax/ebx for invalid subleaves).
+     * (if so, we also get 0 in eax/ebx for invalid subleaves). Zhaoxin implements this too.
      * However AMD rather says that the 0x80000026/0x0b loop should stop when we get "no thread at this level" (0 in ebx[0:15]).
-     * Zhaoxin follows the Intel specs but also returns "no thread at this level" for the last *valid* level (at least on KH-4000).
-     * From the Linux kernel code, it's very likely that AMD also returns "invalid domain"
-     * (because detect_extended_topology() uses that for all x86 CPUs)
-     * but keep with the official doc until AMD can clarify that (see #593).
+     *
+     * Linux kernel <= 6.8 used "invalid domain" for both Intel and AMD (in detect_extended_topology())
+     * but x86 discovery revamp in 6.9 now properly checks both Intel and AMD conditions (in topo_subleaf()).
+     * So let's assume we are allowed to break-out once one of the Intel+AMD conditions is met.
      */
-    if (cpuid_type == amd) {
-      if (!(ebx & 0xffff))
+    if (!(ebx & 0xffff) || !(ecx & 0xff00))
       break;
-    } else {
-      if (!(ecx & 0xff00))
-        break;
-    }
     apic_packageshift = eax & 0x1f;
   }
 
@@ -572,13 +566,8 @@ static void read_extended_topo(struct hwloc_x86_backend_data_s *data, struct pro
 	ecx = level;
 	eax = leaf;
 	cpuid_or_from_dump(&eax, &ebx, &ecx, &edx, src_cpuiddump);
-        if (cpuid_type == amd) {
-          if (!(ebx & 0xffff))
+        if (!(ebx & 0xffff) || !(ecx & 0xff00))
           break;
-        } else {
-          if (!(ecx & 0xff00))
-            break;
-        }
 	apic_nextshift = eax & 0x1f;
 	apic_type = (ecx & 0xff00) >> 8;
 	apic_id = edx;
@@ -1825,7 +1814,7 @@ hwloc_x86_check_cpuiddump_input(const char *src_cpuiddump_path, hwloc_bitmap_t s
     goto out_with_path;
   }
   fclose(file);
-  if (strcmp(line, "Architecture: x86\n")) {
+  if (strncmp(line, "Architecture: x86", 17)) {
     fprintf(stderr, "hwloc/x86: Found non-x86 dumped cpuid summary in %s: %s\n", path, line);
     goto out_with_path;
   }

src/3rdparty/hwloc/src/topology-xml-nolibxml.c

@@ -1,6 +1,6 @@
 /*
  * Copyright © 2009 CNRS
- * Copyright © 2009-2020 Inria.  All rights reserved.
+ * Copyright © 2009-2024 Inria.  All rights reserved.
  * Copyright © 2009-2011 Université Bordeaux
  * Copyright © 2009-2011 Cisco Systems, Inc.  All rights reserved.
  * See COPYING in top-level directory.
@@ -41,7 +41,7 @@ typedef struct hwloc__nolibxml_import_state_data_s {
 static char *
 hwloc__nolibxml_import_ignore_spaces(char *buffer)
 {
-  return buffer + strspn(buffer, " \t\n");
+  return buffer + strspn(buffer, " \t\n\r");
 }
 
 static int

src/3rdparty/hwloc/src/topology-xml.c

@@ -1,6 +1,6 @@
 /*
  * Copyright © 2009 CNRS
- * Copyright © 2009-2023 Inria.  All rights reserved.
+ * Copyright © 2009-2024 Inria.  All rights reserved.
  * Copyright © 2009-2011, 2020 Université Bordeaux
  * Copyright © 2009-2018 Cisco Systems, Inc.  All rights reserved.
  * See COPYING in top-level directory.
@@ -872,6 +872,10 @@ hwloc__xml_import_object(hwloc_topology_t topology,
 	  /* deal with possible future type */
 	  obj->type = HWLOC_OBJ_GROUP;
 	  obj->attr->group.kind = HWLOC_GROUP_KIND_INTEL_MODULE;
+	} else if (!strcasecmp(attrvalue, "Cluster")) {
+	  /* deal with possible future type */
+	  obj->type = HWLOC_OBJ_GROUP;
+	  obj->attr->group.kind = HWLOC_GROUP_KIND_LINUX_CLUSTER;
 	} else if (!strcasecmp(attrvalue, "MemCache")) {
 	  /* ignore possible future type */
 	  obj->type = _HWLOC_OBJ_FUTURE;
@@ -1344,7 +1348,7 @@ hwloc__xml_v2import_support(hwloc_topology_t topology,
     HWLOC_BUILD_ASSERT(sizeof(struct hwloc_topology_support) == 4*sizeof(void*));
     HWLOC_BUILD_ASSERT(sizeof(struct hwloc_topology_discovery_support) == 6);
     HWLOC_BUILD_ASSERT(sizeof(struct hwloc_topology_cpubind_support) == 11);
-    HWLOC_BUILD_ASSERT(sizeof(struct hwloc_topology_membind_support) == 15);
+    HWLOC_BUILD_ASSERT(sizeof(struct hwloc_topology_membind_support) == 16);
     HWLOC_BUILD_ASSERT(sizeof(struct hwloc_topology_misc_support) == 1);
 #endif
 
@@ -1378,6 +1382,7 @@ hwloc__xml_v2import_support(hwloc_topology_t topology,
     else DO(membind,firsttouch_membind);
     else DO(membind,bind_membind);
     else DO(membind,interleave_membind);
+    else DO(membind,weighted_interleave_membind);
     else DO(membind,nexttouch_membind);
     else DO(membind,migrate_membind);
     else DO(membind,get_area_memlocation);
@@ -1436,6 +1441,10 @@ hwloc__xml_v2import_distances(hwloc_topology_t topology,
     }
     else if (!strcmp(attrname, "kind")) {
       kind = strtoul(attrvalue, NULL, 10);
+      /* forward compat with "HOPS" kind in v3 */
+      if (kind & (1UL<<5))
+        /* hops becomes latency */
+        kind = (kind & ~(1UL<<5)) | HWLOC_DISTANCES_KIND_MEANS_LATENCY;
     }
     else if (!strcmp(attrname, "name")) {
       name = attrvalue;
@@ -3087,7 +3096,7 @@ hwloc__xml_v2export_support(hwloc__xml_export_state_t parentstate, hwloc_topolog
   HWLOC_BUILD_ASSERT(sizeof(struct hwloc_topology_support) == 4*sizeof(void*));
   HWLOC_BUILD_ASSERT(sizeof(struct hwloc_topology_discovery_support) == 6);
   HWLOC_BUILD_ASSERT(sizeof(struct hwloc_topology_cpubind_support) == 11);
-  HWLOC_BUILD_ASSERT(sizeof(struct hwloc_topology_membind_support) == 15);
+  HWLOC_BUILD_ASSERT(sizeof(struct hwloc_topology_membind_support) == 16);
   HWLOC_BUILD_ASSERT(sizeof(struct hwloc_topology_misc_support) == 1);
 #endif
 
@@ -3132,6 +3141,7 @@ hwloc__xml_v2export_support(hwloc__xml_export_state_t parentstate, hwloc_topolog
   DO(membind,firsttouch_membind);
   DO(membind,bind_membind);
   DO(membind,interleave_membind);
+  DO(membind,weighted_interleave_membind);
   DO(membind,nexttouch_membind);
   DO(membind,migrate_membind);
   DO(membind,get_area_memlocation);

src/3rdparty/hwloc/src/topology.c

@@ -465,6 +465,20 @@ hwloc_debug_print_objects(int indent __hwloc_attribute_unused, hwloc_obj_t obj)
 #define hwloc_debug_print_objects(indent, obj) do { /* nothing */ } while (0)
 #endif /* !HWLOC_DEBUG */
 
+int hwloc_obj_set_subtype(hwloc_topology_t topology __hwloc_attribute_unused, hwloc_obj_t obj, const char *subtype)
+{
+  char *new = NULL;
+  if (subtype) {
+    new = strdup(subtype);
+    if (!new)
+      return -1;
+  }
+  if (obj->subtype)
+    free(obj->subtype);
+  obj->subtype = new;
+  return 0;
+}
+
 void hwloc__free_infos(struct hwloc_info_s *infos, unsigned count)
 {
   unsigned i;

src/App.cpp

@@ -6,8 +6,8 @@
  * Copyright 2016      Jay D Dee   <jayddee246@gmail.com>
  * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
  * Copyright 2018      Lee Clagett <https://github.com/vtnerd>
- * Copyright 2018-2020 SChernykh   <https://github.com/SChernykh>
- * Copyright 2016-2020 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2024 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
@@ -23,7 +23,6 @@
  *   along with this program. If not, see <http://www.gnu.org/licenses/>.
  */
 
-
 #include <cstdlib>
 #include <uv.h>
 
@@ -61,13 +60,13 @@ int xmrig::App::exec()
         return 2;
     }
 
-    m_signals = std::make_shared<Signals>(this);
-
     int rc = 0;
     if (background(rc)) {
         return rc;
     }
 
+    m_signals = std::make_shared<Signals>(this);
+
     rc = m_controller->init();
     if (rc != 0) {
         return rc;

src/App_unix.cpp

@@ -5,8 +5,8 @@
  * 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-2020 SChernykh   <https://github.com/SChernykh>
- * Copyright 2016-2020 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2024 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
@@ -22,7 +22,6 @@
  *   along with this program. If not, see <http://www.gnu.org/licenses/>.
  */
 
-
 #include <cstdlib>
 #include <csignal>
 #include <cerrno>
@@ -53,16 +52,9 @@ bool xmrig::App::background(int &rc)
         return true;
     }
 
-    i = setsid();
-
-    if (i < 0) {
+    if (setsid() < 0) {
         LOG_ERR("setsid() failed (errno = %d)", errno);
     }
 
-    i = chdir("/");
-    if (i < 0) {
-        LOG_ERR("chdir() failed (errno = %d)", errno);
-    }
-
     return false;
 }

src/backend/common/Hashrate.cpp

@@ -30,10 +30,10 @@
 #include "base/tools/Handle.h"
 
 
-inline static const char *format(double h, char *buf, size_t size)
+inline static const char *format(std::pair<bool, double> h, char *buf, size_t size)
 {
-    if (std::isnormal(h)) {
-        snprintf(buf, size, (h < 100.0) ? "%04.2f" : "%03.1f", h);
+    if (h.first) {
+        snprintf(buf, size, (h.second < 100.0) ? "%04.2f" : "%03.1f", h.second);
         return buf;
     }
 
@@ -80,15 +80,16 @@ double xmrig::Hashrate::average() const
 }
 
 
-const char *xmrig::Hashrate::format(double h, char *buf, size_t size)
+const char *xmrig::Hashrate::format(std::pair<bool, double> h, char *buf, size_t size)
 {
     return ::format(h, buf, size);
 }
 
 
-rapidjson::Value xmrig::Hashrate::normalize(double d)
+rapidjson::Value xmrig::Hashrate::normalize(std::pair<bool, double> d)
 {
-    return Json::normalize(d, false);
+    using namespace rapidjson;
+    return d.first ? Value(floor(d.second * 100.0) / 100.0) : Value(kNullType);
 }
 
 
@@ -122,11 +123,11 @@ rapidjson::Value xmrig::Hashrate::toJSON(size_t threadId, rapidjson::Document &d
 #endif
 
 
-double xmrig::Hashrate::hashrate(size_t index, size_t ms) const
+std::pair<bool, double> xmrig::Hashrate::hashrate(size_t index, size_t ms) const
 {
     assert(index < m_threads);
     if (index >= m_threads) {
-        return nan("");
+        return { false, 0.0 };
     }
 
     uint64_t earliestHashCount = 0;
@@ -157,17 +158,27 @@ double xmrig::Hashrate::hashrate(size_t index, size_t ms) const
     } while (idx != idx_start);
 
     if (!haveFullSet || earliestStamp == 0 || lastestStamp == 0) {
-        return nan("");
+        return { false, 0.0 };
     }
 
-    if (lastestStamp - earliestStamp == 0) {
-        return nan("");
+    if (lastestHashCnt == earliestHashCount) {
+        return { true, 0.0 };
+    }
+
+    if (lastestStamp == earliestStamp) {
+        return { false, 0.0 };
     }
 
     const auto hashes = static_cast<double>(lastestHashCnt - earliestHashCount);
-    const auto time   = static_cast<double>(lastestStamp - earliestStamp) / 1000.0;
+    const auto time   = static_cast<double>(lastestStamp - earliestStamp);
 
-    return hashes / time;
+    const auto hr = hashes * 1000.0 / time;
+
+    if (!std::isnormal(hr)) {
+        return { false, 0.0 };
+    }
+
+    return { true, hr };
 }
 
 

src/backend/common/Hashrate.h

@@ -47,16 +47,16 @@ public:
     Hashrate(size_t threads);
     ~Hashrate();
 
-    inline double calc(size_t ms) const                                     { const double data = hashrate(0U, ms); return std::isnormal(data) ? data : 0.0; }
-    inline double calc(size_t threadId, size_t ms) const                    { return hashrate(threadId + 1, ms); }
+    inline std::pair<bool, double> calc(size_t ms) const                    { return hashrate(0U, ms); }
+    inline std::pair<bool, double> calc(size_t threadId, size_t ms) const   { return hashrate(threadId + 1, ms); }
     inline size_t threads() const                                           { return m_threads > 0U ? m_threads - 1U : 0U; }
     inline void add(size_t threadId, uint64_t count, uint64_t timestamp)    { addData(threadId + 1U, count, timestamp); }
     inline void add(uint64_t count, uint64_t timestamp)                     { addData(0U, count, timestamp); }
 
     double average() const;
 
-    static const char *format(double h, char *buf, size_t size);
-    static rapidjson::Value normalize(double d);
+    static const char *format(std::pair<bool, double> h, char *buf, size_t size);
+    static rapidjson::Value normalize(std::pair<bool, double> d);
 
 #   ifdef XMRIG_FEATURE_API
     rapidjson::Value toJSON(rapidjson::Document &doc) const;
@@ -64,7 +64,7 @@ public:
 #   endif
 
 private:
-    double hashrate(size_t index, size_t ms) const;
+    std::pair<bool, double> hashrate(size_t index, size_t ms) const;
     void addData(size_t index, uint64_t count, uint64_t timestamp);
 
     constexpr static size_t kBucketSize = 2 << 11;

src/backend/common/Thread.h

@@ -65,22 +65,22 @@ public:
         }
     }
 #   else
-    inline ~Thread() { m_thread.join(); delete m_worker; }
+    inline ~Thread() { m_thread.join(); }
 
     inline void start(void *(*callback)(void *))    { m_thread = std::thread(callback, this); }
 #   endif
 
     inline const T &config() const                  { return m_config; }
     inline IBackend *backend() const                { return m_backend; }
-    inline IWorker *worker() const                  { return m_worker; }
+    inline IWorker* worker() const                  { return m_worker.get(); }
     inline size_t id() const                        { return m_id; }
-    inline void setWorker(IWorker *worker)          { m_worker = worker; }
+    inline void setWorker(std::shared_ptr<IWorker> worker) { m_worker = worker; }
 
 private:
     const size_t m_id    = 0;
     const T m_config;
     IBackend *m_backend;
-    IWorker *m_worker       = nullptr;
+    std::shared_ptr<IWorker> m_worker;
 
     #ifdef XMRIG_OS_APPLE
     pthread_t m_thread{};

src/backend/common/Workers.cpp

@@ -62,19 +62,12 @@ public:
 
 template<class T>
 xmrig::Workers<T>::Workers() :
-    d_ptr(new WorkersPrivate())
+    d_ptr(std::make_shared<WorkersPrivate>())
 {
 
 }
 
 
-template<class T>
-xmrig::Workers<T>::~Workers()
-{
-    delete d_ptr;
-}
-
-
 template<class T>
 bool xmrig::Workers<T>::tick(uint64_t)
 {
@@ -88,7 +81,7 @@ bool xmrig::Workers<T>::tick(uint64_t)
     uint64_t hashCount      = 0;
     uint64_t rawHashes      = 0;
 
-    for (Thread<T> *handle : m_workers) {
+    for (auto& handle : m_workers) {
         IWorker *worker = handle->worker();
         if (worker) {
             worker->hashrateData(hashCount, ts, rawHashes);
@@ -135,10 +128,6 @@ void xmrig::Workers<T>::stop()
     Nonce::stop(T::backend());
 #   endif
 
-    for (Thread<T> *worker : m_workers) {
-        delete worker;
-    }
-
     m_workers.clear();
 
 #   ifdef XMRIG_MINER_PROJECT
@@ -166,7 +155,7 @@ void xmrig::Workers<T>::start(const std::vector<T> &data, const std::shared_ptr<
 
 
 template<class T>
-xmrig::IWorker *xmrig::Workers<T>::create(Thread<T> *)
+std::shared_ptr<xmrig::IWorker> xmrig::Workers<T>::create(Thread<T> *)
 {
     return nullptr;
 }
@@ -177,22 +166,21 @@ void *xmrig::Workers<T>::onReady(void *arg)
 {
     auto handle = static_cast<Thread<T>* >(arg);
 
-    IWorker *worker = create(handle);
-    assert(worker != nullptr);
+    std::shared_ptr<IWorker> worker = create(handle);
+    assert(worker);
 
     if (!worker || !worker->selfTest()) {
         LOG_ERR("%s " RED("thread ") RED_BOLD("#%zu") RED(" self-test failed"), T::tag(), worker ? worker->id() : 0);
 
-        handle->backend()->start(worker, false);
-        delete worker;
-
+        worker.reset();
+        handle->backend()->start(worker.get(), false);
         return nullptr;
     }
 
     assert(handle->backend() != nullptr);
 
     handle->setWorker(worker);
-    handle->backend()->start(worker, true);
+    handle->backend()->start(worker.get(), true);
 
     return nullptr;
 }
@@ -202,7 +190,7 @@ template<class T>
 void xmrig::Workers<T>::start(const std::vector<T> &data, bool /*sleep*/)
 {
     for (const auto &item : data) {
-        m_workers.push_back(new Thread<T>(d_ptr->backend, m_workers.size(), item));
+        m_workers.emplace_back(std::make_shared<Thread<T>>(d_ptr->backend, m_workers.size(), item));
     }
 
     d_ptr->hashrate = std::make_shared<Hashrate>(m_workers.size());
@@ -211,7 +199,7 @@ void xmrig::Workers<T>::start(const std::vector<T> &data, bool /*sleep*/)
     Nonce::touch(T::backend());
 #   endif
 
-    for (auto worker : m_workers) {
+    for (auto& worker : m_workers) {
         worker->start(Workers<T>::onReady);
     }
 }
@@ -221,34 +209,34 @@ namespace xmrig {
 
 
 template<>
-xmrig::IWorker *xmrig::Workers<CpuLaunchData>::create(Thread<CpuLaunchData> *handle)
+std::shared_ptr<xmrig::IWorker> Workers<CpuLaunchData>::create(Thread<CpuLaunchData> *handle)
 {
 #   ifdef XMRIG_MINER_PROJECT
     switch (handle->config().intensity) {
     case 1:
-        return new CpuWorker<1>(handle->id(), handle->config());
+        return std::make_shared<CpuWorker<1>>(handle->id(), handle->config());
 
     case 2:
-        return new CpuWorker<2>(handle->id(), handle->config());
+        return std::make_shared<CpuWorker<2>>(handle->id(), handle->config());
 
     case 3:
-        return new CpuWorker<3>(handle->id(), handle->config());
+        return std::make_shared<CpuWorker<3>>(handle->id(), handle->config());
 
     case 4:
-        return new CpuWorker<4>(handle->id(), handle->config());
+        return std::make_shared<CpuWorker<4>>(handle->id(), handle->config());
 
     case 5:
-        return new CpuWorker<5>(handle->id(), handle->config());
+        return std::make_shared<CpuWorker<5>>(handle->id(), handle->config());
 
     case 8:
-        return new CpuWorker<8>(handle->id(), handle->config());
+        return std::make_shared<CpuWorker<8>>(handle->id(), handle->config());
     }
 
     return nullptr;
 #   else
     assert(handle->config().intensity == 1);
 
-    return new CpuWorker<1>(handle->id(), handle->config());
+    return std::make_shared<CpuWorker<1>>(handle->id(), handle->config());
 #   endif
 }
 
@@ -258,9 +246,9 @@ template class Workers<CpuLaunchData>;
 
 #ifdef XMRIG_FEATURE_OPENCL
 template<>
-xmrig::IWorker *xmrig::Workers<OclLaunchData>::create(Thread<OclLaunchData> *handle)
+std::shared_ptr<xmrig::IWorker> Workers<OclLaunchData>::create(Thread<OclLaunchData> *handle)
 {
-    return new OclWorker(handle->id(), handle->config());
+    return std::make_shared<OclWorker>(handle->id(), handle->config());
 }
 
 
@@ -270,9 +258,9 @@ template class Workers<OclLaunchData>;
 
 #ifdef XMRIG_FEATURE_CUDA
 template<>
-xmrig::IWorker *xmrig::Workers<CudaLaunchData>::create(Thread<CudaLaunchData> *handle)
+std::shared_ptr<xmrig::IWorker> Workers<CudaLaunchData>::create(Thread<CudaLaunchData> *handle)
 {
-    return new CudaWorker(handle->id(), handle->config());
+    return std::make_shared<CudaWorker>(handle->id(), handle->config());
 }
 
 

src/backend/common/Workers.h

@@ -52,7 +52,6 @@ public:
     XMRIG_DISABLE_COPY_MOVE(Workers)
 
     Workers();
-    ~Workers();
 
     inline void start(const std::vector<T> &data)   { start(data, true); }
 
@@ -67,20 +66,20 @@ public:
 #   endif
 
 private:
-    static IWorker *create(Thread<T> *handle);
+    static std::shared_ptr<IWorker> create(Thread<T> *handle);
     static void *onReady(void *arg);
 
     void start(const std::vector<T> &data, bool sleep);
 
-    std::vector<Thread<T> *> m_workers;
-    WorkersPrivate *d_ptr;
+    std::vector<std::shared_ptr<Thread<T>>> m_workers;
+    std::shared_ptr<WorkersPrivate> d_ptr;
 };
 
 
 template<class T>
 void xmrig::Workers<T>::jobEarlyNotification(const Job &job)
 {
-    for (Thread<T>* t : m_workers) {
+    for (auto& t : m_workers) {
         if (t->worker()) {
             t->worker()->jobEarlyNotification(job);
         }
@@ -89,20 +88,20 @@ void xmrig::Workers<T>::jobEarlyNotification(const Job &job)
 
 
 template<>
-IWorker *Workers<CpuLaunchData>::create(Thread<CpuLaunchData> *handle);
+std::shared_ptr<IWorker> Workers<CpuLaunchData>::create(Thread<CpuLaunchData> *handle);
 extern template class Workers<CpuLaunchData>;
 
 
 #ifdef XMRIG_FEATURE_OPENCL
 template<>
-IWorker *Workers<OclLaunchData>::create(Thread<OclLaunchData> *handle);
+std::shared_ptr<IWorker> Workers<OclLaunchData>::create(Thread<OclLaunchData> *handle);
 extern template class Workers<OclLaunchData>;
 #endif
 
 
 #ifdef XMRIG_FEATURE_CUDA
 template<>
-IWorker *Workers<CudaLaunchData>::create(Thread<CudaLaunchData> *handle);
+std::shared_ptr<IWorker> Workers<CudaLaunchData>::create(Thread<CudaLaunchData> *handle);
 extern template class Workers<CudaLaunchData>;
 #endif
 

src/backend/common/benchmark/BenchState.cpp

@@ -51,7 +51,7 @@ public:
 };
 
 
-static BenchStatePrivate *d_ptr = nullptr;
+static std::shared_ptr<BenchStatePrivate> d_ptr;
 std::atomic<uint64_t> BenchState::m_data{};
 
 
@@ -61,7 +61,7 @@ std::atomic<uint64_t> BenchState::m_data{};
 
 bool xmrig::BenchState::isDone()
 {
-    return d_ptr == nullptr;
+    return !d_ptr;
 }
 
 
@@ -105,14 +105,13 @@ uint64_t xmrig::BenchState::start(size_t threads, const IBackend *backend)
 
 void xmrig::BenchState::destroy()
 {
-    delete d_ptr;
-    d_ptr = nullptr;
+    d_ptr.reset();
 }
 
 
 void xmrig::BenchState::done()
 {
-    assert(d_ptr != nullptr && d_ptr->async && d_ptr->remaining > 0);
+    assert(d_ptr && d_ptr->async && d_ptr->remaining > 0);
 
     const uint64_t ts = Chrono::steadyMSecs();
 
@@ -129,15 +128,15 @@ void xmrig::BenchState::done()
 
 void xmrig::BenchState::init(IBenchListener *listener, uint32_t size)
 {
-    assert(d_ptr == nullptr);
+    assert(!d_ptr);
 
-    d_ptr = new BenchStatePrivate(listener, size);
+    d_ptr = std::make_shared<BenchStatePrivate>(listener, size);
 }
 
 
 void xmrig::BenchState::setSize(uint32_t size)
 {
-    assert(d_ptr != nullptr);
+    assert(d_ptr);
 
     d_ptr->size = size;
 }

src/backend/common/misc/PciTopology.h

@@ -1,6 +1,6 @@
 /* XMRig
- * Copyright (c) 2018-2021 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2021 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright (c) 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright (c) 2016-2024 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
@@ -19,10 +19,8 @@
 #ifndef XMRIG_PCITOPOLOGY_H
 #define XMRIG_PCITOPOLOGY_H
 
-
 #include <cstdio>
 
-
 #include "base/tools/String.h"
 
 
@@ -33,7 +31,14 @@ class PciTopology
 {
 public:
     PciTopology() = default;
-    PciTopology(uint32_t bus, uint32_t device, uint32_t function) : m_valid(true), m_bus(bus), m_device(device), m_function(function) {}
+
+    template<typename T>
+    inline PciTopology(T bus, T device, T function)
+        : m_valid(true),
+          m_bus(static_cast<uint8_t>(bus)),
+          m_device(static_cast<uint8_t>(device)),
+          m_function(static_cast<uint8_t>(function))
+    {}
 
     inline bool isEqual(const PciTopology &other) const     { return m_valid == other.m_valid && toUint32() == other.toUint32(); }
     inline bool isValid() const                             { return m_valid; }
@@ -70,4 +75,4 @@ private:
 } // namespace xmrig
 
 
-#endif /* XMRIG_PCITOPOLOGY_H */
+#endif // XMRIG_PCITOPOLOGY_H

src/backend/cpu/Cpu.cpp

@@ -31,20 +31,20 @@
 #endif
 
 
-static xmrig::ICpuInfo *cpuInfo = nullptr;
+static std::shared_ptr<xmrig::ICpuInfo> cpuInfo;
 
 
 xmrig::ICpuInfo *xmrig::Cpu::info()
 {
-    if (cpuInfo == nullptr) {
+    if (!cpuInfo) {
 #       if defined(XMRIG_FEATURE_HWLOC)
-        cpuInfo = new HwlocCpuInfo();
+        cpuInfo = std::make_shared<HwlocCpuInfo>();
 #       else
-        cpuInfo = new BasicCpuInfo();
+        cpuInfo = std::make_shared<BasicCpuInfo>();
 #       endif
     }
 
-    return cpuInfo;
+    return cpuInfo.get();
 }
 
 
@@ -56,6 +56,5 @@ rapidjson::Value xmrig::Cpu::toJSON(rapidjson::Document &doc)
 
 void xmrig::Cpu::release()
 {
-    delete cpuInfo;
-    cpuInfo = nullptr;
+    cpuInfo.reset();
 }

src/backend/cpu/CpuBackend.cpp

@@ -242,7 +242,7 @@ const char *xmrig::cpu_tag()
 
 
 xmrig::CpuBackend::CpuBackend(Controller *controller) :
-    d_ptr(new CpuBackendPrivate(controller))
+    d_ptr(std::make_shared<CpuBackendPrivate>(controller))
 {
     d_ptr->workers.setBackend(this);
 }
@@ -250,7 +250,6 @@ xmrig::CpuBackend::CpuBackend(Controller *controller) :
 
 xmrig::CpuBackend::~CpuBackend()
 {
-    delete d_ptr;
 }
 
 

src/backend/cpu/CpuBackend.h

@@ -70,7 +70,7 @@ protected:
 #   endif
 
 private:
-    CpuBackendPrivate *d_ptr;
+    std::shared_ptr<CpuBackendPrivate> d_ptr;
 };
 
 

src/backend/cpu/CpuConfig_gen.h

@@ -1,6 +1,6 @@
 /* XMRig
- * Copyright (c) 2018-2021 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2021 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright (c) 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright (c) 2016-2024 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
@@ -122,17 +122,6 @@ size_t inline generate<Algorithm::RANDOM_X>(Threads<CpuThreads> &threads, uint32
         }
     }
 
-    if (!threads.isExist(Algorithm::RX_KEVA)) {
-        auto keva = cpuInfo->threads(Algorithm::RX_KEVA, limit);
-        if (keva == wow) {
-            threads.setAlias(Algorithm::RX_KEVA, Algorithm::kRX_WOW);
-            ++count;
-        }
-        else {
-            count += threads.move(Algorithm::kRX_KEVA, std::move(keva));
-        }
-    }
-
     if (!threads.isExist(Algorithm::RX_WOW)) {
         count += threads.move(Algorithm::kRX_WOW, std::move(wow));
     }

src/backend/cpu/CpuWorker.cpp

@@ -57,7 +57,7 @@ static constexpr uint32_t kReserveCount = 32768;
 
 #ifdef XMRIG_ALGO_CN_HEAVY
 static std::mutex cn_heavyZen3MemoryMutex;
-VirtualMemory* cn_heavyZen3Memory = nullptr;
+std::shared_ptr<VirtualMemory> cn_heavyZen3Memory;
 #endif
 
 } // namespace xmrig
@@ -87,14 +87,14 @@ xmrig::CpuWorker<N>::CpuWorker(size_t id, const CpuLaunchData &data) :
         if (!cn_heavyZen3Memory) {
             // Round up number of threads to the multiple of 8
             const size_t num_threads = ((m_threads + 7) / 8) * 8;
-            cn_heavyZen3Memory = new VirtualMemory(m_algorithm.l3() * num_threads, data.hugePages, false, false, node());
+            cn_heavyZen3Memory = std::make_shared<VirtualMemory>(m_algorithm.l3() * num_threads, data.hugePages, false, false, node());
         }
         m_memory = cn_heavyZen3Memory;
     }
     else
 #   endif
     {
-        m_memory = new VirtualMemory(m_algorithm.l3() * N, data.hugePages, false, true, node());
+        m_memory = std::make_shared<VirtualMemory>(m_algorithm.l3() * N, data.hugePages, false, true, node());
     }
 
 #   ifdef XMRIG_ALGO_GHOSTRIDER
@@ -107,7 +107,7 @@ template<size_t N>
 xmrig::CpuWorker<N>::~CpuWorker()
 {
 #   ifdef XMRIG_ALGO_RANDOMX
-    RxVm::destroy(m_vm);
+    m_vm.reset();
 #   endif
 
     CnCtx::release(m_ctx, N);
@@ -116,7 +116,7 @@ xmrig::CpuWorker<N>::~CpuWorker()
     if (m_memory != cn_heavyZen3Memory)
 #   endif
     {
-        delete m_memory;
+        m_memory.reset();
     }
 
 #   ifdef XMRIG_ALGO_GHOSTRIDER
@@ -148,7 +148,7 @@ void xmrig::CpuWorker<N>::allocateRandomX_VM()
     }
     else if (!dataset->get() && (m_job.currentJob().seed() != m_seed)) {
         // Update RandomX light VM with the new seed
-        randomx_vm_set_cache(m_vm, dataset->cache()->get());
+        randomx_vm_set_cache(m_vm.get(), dataset->cache()->get());
     }
     m_seed = m_job.currentJob().seed();
 }
@@ -296,7 +296,7 @@ void xmrig::CpuWorker<N>::start()
                     if (job.hasMinerSignature()) {
                         job.generateMinerSignature(m_job.blob(), job.size(), miner_signature_ptr);
                     }
-                    randomx_calculate_hash_first(m_vm, tempHash, m_job.blob(), job.size());
+                    randomx_calculate_hash_first(m_vm.get(), tempHash, m_job.blob(), job.size());
                 }
 
                 if (!nextRound()) {
@@ -307,7 +307,7 @@ void xmrig::CpuWorker<N>::start()
                     memcpy(miner_signature_saved, miner_signature_ptr, sizeof(miner_signature_saved));
                     job.generateMinerSignature(m_job.blob(), job.size(), miner_signature_ptr);
                 }
-                randomx_calculate_hash_next(m_vm, tempHash, m_job.blob(), job.size(), m_hash);
+                randomx_calculate_hash_next(m_vm.get(), tempHash, m_job.blob(), job.size(), m_hash);
             }
             else
 #           endif
@@ -359,9 +359,11 @@ void xmrig::CpuWorker<N>::start()
             }
         }
 
+        if (!Nonce::isPaused()) {
             consumeJob();
         }
     }
+}
 
 
 template<size_t N>

src/backend/cpu/CpuWorker.h

@@ -66,7 +66,7 @@ protected:
     void hashrateData(uint64_t &hashCount, uint64_t &timeStamp, uint64_t &rawHashes) const override;
     void start() override;
 
-    inline const VirtualMemory *memory() const override     { return m_memory; }
+    inline const VirtualMemory* memory() const override     { return m_memory.get(); }
     inline size_t intensity() const override                { return N; }
     inline void jobEarlyNotification(const Job&) override   {}
 
@@ -92,11 +92,11 @@ private:
     const Miner *m_miner;
     const size_t m_threads;
     cryptonight_ctx *m_ctx[N];
-    VirtualMemory *m_memory = nullptr;
+    std::shared_ptr<VirtualMemory> m_memory;
     WorkerJob<N> m_job;
 
 #   ifdef XMRIG_ALGO_RANDOMX
-    randomx_vm *m_vm        = nullptr;
+    std::shared_ptr<randomx_vm> m_vm;
     Buffer m_seed;
 #   endif
 

src/backend/cpu/interfaces/ICpuInfo.h

@@ -52,7 +52,8 @@ public:
         ARCH_ZEN_PLUS,
         ARCH_ZEN2,
         ARCH_ZEN3,
-        ARCH_ZEN4
+        ARCH_ZEN4,
+        ARCH_ZEN5
     };
 
     enum MsrMod : uint32_t {
@@ -60,12 +61,13 @@ public:
         MSR_MOD_RYZEN_17H,
         MSR_MOD_RYZEN_19H,
         MSR_MOD_RYZEN_19H_ZEN4,
+        MSR_MOD_RYZEN_1AH_ZEN5,
         MSR_MOD_INTEL,
         MSR_MOD_CUSTOM,
         MSR_MOD_MAX
     };
 
-#   define MSR_NAMES_LIST "none", "ryzen_17h", "ryzen_19h", "ryzen_19h_zen4", "intel", "custom"
+#   define MSR_NAMES_LIST "none", "ryzen_17h", "ryzen_19h", "ryzen_19h_zen4", "ryzen_1Ah_zen5", "intel", "custom"
 
     enum Flag : uint32_t {
         FLAG_AES,

src/backend/cpu/platform/BasicCpuInfo.cpp

@@ -64,7 +64,7 @@ static_assert(kCpuFlagsSize == ICpuInfo::FLAG_MAX, "kCpuFlagsSize and FLAG_MAX m
 
 
 #ifdef XMRIG_FEATURE_MSR
-constexpr size_t kMsrArraySize                                  = 6;
+constexpr size_t kMsrArraySize                                  = 7;
 static const std::array<const char *, kMsrArraySize> msrNames   = { MSR_NAMES_LIST };
 static_assert(kMsrArraySize == ICpuInfo::MSR_MOD_MAX, "kMsrArraySize and MSR_MOD_MAX mismatch");
 #endif
@@ -260,6 +260,11 @@ xmrig::BasicCpuInfo::BasicCpuInfo() :
                     }
                     break;
 
+                case 0x1a:
+                    m_arch = ARCH_ZEN5;
+                    m_msrMod = MSR_MOD_RYZEN_1AH_ZEN5;
+                    break;
+
                 default:
                     m_msrMod = MSR_MOD_NONE;
                     break;

src/backend/cpu/platform/HwlocCpuInfo.cpp

@@ -326,7 +326,8 @@ void xmrig::HwlocCpuInfo::processTopLevelCache(hwloc_obj_t cache, const Algorith
         }
     }
 
-    if (scratchpad == 2 * oneMiB) {
+    // This code is supposed to run only on Intel CPUs
+    if ((vendor() == VENDOR_INTEL) && (scratchpad == 2 * oneMiB)) {
         if (L2 && (cores.size() * oneMiB) == L2 && L2_associativity == 16 && L3 >= L2) {
             L3    = L2;
             extra = L2;
@@ -341,7 +342,7 @@ void xmrig::HwlocCpuInfo::processTopLevelCache(hwloc_obj_t cache, const Algorith
     }
 
 #   ifdef XMRIG_ALGO_RANDOMX
-    if ((algorithm.family() == Algorithm::RANDOM_X) && L3_exclusive && (PUs > cores.size()) && (PUs < cores.size() * 2)) {
+    if ((vendor() == VENDOR_INTEL) && (algorithm.family() == Algorithm::RANDOM_X) && L3_exclusive && (PUs < cores.size() * 2)) {
         // Use all L3+L2 on latest Intel CPUs with P-cores, E-cores and exclusive L3 cache
         cacheHashes = (L3 + L2) / scratchpad;
     }

src/backend/cuda/CudaBackend.cpp

@@ -1,6 +1,6 @@
 /* XMRig
- * Copyright (c) 2018-2021 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2021 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright (c) 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright (c) 2016-2024 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
@@ -227,7 +227,7 @@ public:
 #           endif
 
             Log::print("|" CYAN_BOLD("%3zu") " |" CYAN_BOLD("%4u") " |" YELLOW(" %7s") " |" CYAN_BOLD("%10d") " |" CYAN_BOLD("%8d") " |"
-                       CYAN_BOLD("%7d") " |" CYAN_BOLD("%3d") " |" CYAN_BOLD("%4d") " |" CYAN("%7zu") " | " GREEN("%s"),
+                       CYAN_BOLD("%7d") " |" CYAN_BOLD("%3d") " |" CYAN_BOLD("%4d") " |" CYAN("%7zu") " | " GREEN_BOLD("%s"),
                        i,
                        data.thread.index(),
                        data.device.topology().toString().data(),
@@ -372,15 +372,20 @@ void xmrig::CudaBackend::printHashrate(bool details)
 
     char num[16 * 3] = { 0 };
 
-    const double hashrate_short  = hashrate()->calc(Hashrate::ShortInterval);
-    const double hashrate_medium = hashrate()->calc(Hashrate::MediumInterval);
-    const double hashrate_large  = hashrate()->calc(Hashrate::LargeInterval);
+    auto hashrate_short  = hashrate()->calc(Hashrate::ShortInterval);
+    auto hashrate_medium = hashrate()->calc(Hashrate::MediumInterval);
+    auto hashrate_large  = hashrate()->calc(Hashrate::LargeInterval);
 
     double scale = 1.0;
     const char* h = " H/s";
 
-    if ((hashrate_short >= 1e6) || (hashrate_medium >= 1e6) || (hashrate_large >= 1e6)) {
+    if ((hashrate_short.second >= 1e6) || (hashrate_medium.second >= 1e6) || (hashrate_large.second >= 1e6)) {
         scale = 1e-6;
+
+        hashrate_short.second  *= scale;
+        hashrate_medium.second *= scale;
+        hashrate_large.second  *= scale;
+
         h = "MH/s";
     }
 
@@ -388,12 +393,20 @@ void xmrig::CudaBackend::printHashrate(bool details)
 
     size_t i = 0;
     for (const auto& data : d_ptr->threads) {
+        auto h0 = hashrate()->calc(i, Hashrate::ShortInterval);
+        auto h1 = hashrate()->calc(i, Hashrate::MediumInterval);
+        auto h2 = hashrate()->calc(i, Hashrate::LargeInterval);
+
+        h0.second *= scale;
+        h1.second *= scale;
+        h2.second *= scale;
+
         Log::print("| %8zu | %8" PRId64 " | %8s | %8s | %8s |" CYAN_BOLD(" #%u") YELLOW(" %s") GREEN(" %s"),
                     i,
                     data.thread.affinity(),
-                    Hashrate::format(hashrate()->calc(i, Hashrate::ShortInterval)  * scale, num,          sizeof num / 3),
-                    Hashrate::format(hashrate()->calc(i, Hashrate::MediumInterval) * scale, num + 16,      sizeof num / 3),
-                    Hashrate::format(hashrate()->calc(i, Hashrate::LargeInterval)  * scale, num + 16 * 2, sizeof num / 3),
+                    Hashrate::format(h0, num,          sizeof num / 3),
+                    Hashrate::format(h1, num + 16,     sizeof num / 3),
+                    Hashrate::format(h2, num + 16 * 2, sizeof num / 3),
                     data.device.index(),
                     data.device.topology().toString().data(),
                     data.device.name().data()
@@ -403,9 +416,9 @@ void xmrig::CudaBackend::printHashrate(bool details)
     }
 
     Log::print(WHITE_BOLD_S "|        - |        - | %8s | %8s | %8s |",
-               Hashrate::format(hashrate_short  * scale, num,          sizeof num / 3),
-               Hashrate::format(hashrate_medium * scale, num + 16,     sizeof num / 3),
-               Hashrate::format(hashrate_large  * scale, num + 16 * 2, sizeof num / 3)
+               Hashrate::format(hashrate_short , num,          sizeof num / 3),
+               Hashrate::format(hashrate_medium, num + 16,     sizeof num / 3),
+               Hashrate::format(hashrate_large , num + 16 * 2, sizeof num / 3)
                );
 }
 

src/backend/cuda/CudaConfig_gen.h

@@ -114,7 +114,6 @@ size_t inline generate<Algorithm::RANDOM_X>(Threads<CudaThreads> &threads, const
     auto rx  = CudaThreads(devices, Algorithm::RX_0);
     auto wow = CudaThreads(devices, Algorithm::RX_WOW);
     auto arq = CudaThreads(devices, Algorithm::RX_ARQ);
-    auto kva = CudaThreads(devices, Algorithm::RX_KEVA);
 
     if (!threads.isExist(Algorithm::RX_WOW) && wow != rx) {
         count += threads.move(Algorithm::kRX_WOW, std::move(wow));
@@ -124,10 +123,6 @@ size_t inline generate<Algorithm::RANDOM_X>(Threads<CudaThreads> &threads, const
         count += threads.move(Algorithm::kRX_ARQ, std::move(arq));
     }
 
-    if (!threads.isExist(Algorithm::RX_KEVA) && kva != rx) {
-        count += threads.move(Algorithm::kRX_KEVA, std::move(kva));
-    }
-
     count += threads.move(Algorithm::kRX, std::move(rx));
 
     return count;

src/backend/cuda/CudaWorker.cpp

@@ -158,7 +158,7 @@ void xmrig::CudaWorker::start()
             std::this_thread::yield();
         }
 
-        if (!consumeJob()) {
+        if (isReady() && !consumeJob()) {
             return;
         }
     }

src/backend/cuda/runners/CudaRxRunner.cpp

@@ -5,8 +5,8 @@
  * 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>
+ * Copyright 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2024 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
@@ -22,7 +22,6 @@
  *   along with this program. If not, see <http://www.gnu.org/licenses/>.
  */
 
-
 #include "backend/cuda/runners/CudaRxRunner.h"
 #include "backend/cuda/CudaLaunchData.h"
 #include "backend/cuda/wrappers/CudaLib.h"
@@ -55,12 +54,21 @@ bool xmrig::CudaRxRunner::run(uint32_t startNonce, uint32_t *rescount, uint32_t
 
 bool xmrig::CudaRxRunner::set(const Job &job, uint8_t *blob)
 {
+    if (!m_datasetHost && (m_seed != job.seed())) {
+        m_seed = job.seed();
+
+        if (m_ready) {
+            const auto *dataset = Rx::dataset(job, 0);
+            callWrapper(CudaLib::rxUpdateDataset(m_ctx, dataset->raw(), dataset->size(false)));
+        }
+    }
+
     const bool rc = CudaBaseRunner::set(job, blob);
     if (!rc || m_ready) {
         return rc;
     }
 
-    auto dataset = Rx::dataset(job, 0);
+    const auto *dataset = Rx::dataset(job, 0);
     m_ready = callWrapper(CudaLib::rxPrepare(m_ctx, dataset->raw(), dataset->size(false), m_datasetHost, m_intensity));
 
     return m_ready;

src/backend/cuda/runners/CudaRxRunner.h

@@ -5,8 +5,8 @@
  * 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>
+ * Copyright 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2024 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
@@ -27,6 +27,7 @@
 
 
 #include "backend/cuda/runners/CudaBaseRunner.h"
+#include "base/tools/Buffer.h"
 
 
 namespace xmrig {
@@ -46,6 +47,7 @@ protected:
 private:
     bool m_ready             = false;
     const bool m_datasetHost = false;
+    Buffer m_seed;
     size_t m_intensity       = 0;
 };
 

src/backend/cuda/wrappers/CudaDevice.cpp

@@ -5,8 +5,8 @@
  * 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-2020 SChernykh   <https://github.com/SChernykh>
- * Copyright 2016-2020 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2024 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
@@ -22,7 +22,6 @@
  *   along with this program. If not, see <http://www.gnu.org/licenses/>.
  */
 
-
 #include "backend/cuda/wrappers/CudaDevice.h"
 #include "3rdparty/rapidjson/document.h"
 #include "backend/cuda/CudaThreads.h"
@@ -41,7 +40,7 @@
 xmrig::CudaDevice::CudaDevice(uint32_t index, int32_t bfactor, int32_t bsleep) :
     m_index(index)
 {
-    auto ctx = CudaLib::alloc(index, bfactor, bsleep);
+    auto *ctx = CudaLib::alloc(index, bfactor, bsleep);
     if (!CudaLib::deviceInfo(ctx, 0, 0, Algorithm::INVALID)) {
         CudaLib::release(ctx);
 
@@ -50,7 +49,7 @@ xmrig::CudaDevice::CudaDevice(uint32_t index, int32_t bfactor, int32_t bsleep) :
 
     m_ctx       = ctx;
     m_name      = CudaLib::deviceName(ctx);
-    m_topology  = PciTopology(CudaLib::deviceUint(ctx, CudaLib::DevicePciBusID), CudaLib::deviceUint(ctx, CudaLib::DevicePciDeviceID), 0);
+    m_topology  = { CudaLib::deviceUint(ctx, CudaLib::DevicePciBusID), CudaLib::deviceUint(ctx, CudaLib::DevicePciDeviceID), 0U };
 }
 
 

src/backend/cuda/wrappers/CudaLib.cpp

@@ -19,10 +19,10 @@
 #include <stdexcept>
 #include <uv.h>
 
-
 #include "backend/cuda/wrappers/CudaLib.h"
 #include "base/io/Env.h"
 #include "base/io/log/Log.h"
+#include "base/io/log/Tags.h"
 #include "base/kernel/Process.h"
 #include "crypto/rx/RxAlgo.h"
 
@@ -68,6 +68,7 @@ static const char *kPluginVersion                       = "pluginVersion";
 static const char *kRelease                             = "release";
 static const char *kRxHash                              = "rxHash";
 static const char *kRxPrepare                           = "rxPrepare";
+static const char *kRxUpdateDataset                     = "rxUpdateDataset";
 static const char *kSetJob                              = "setJob";
 static const char *kSetJob_v2                           = "setJob_v2";
 static const char *kVersion                             = "version";
@@ -92,6 +93,7 @@ using pluginVersion_t                                   = const char * (*)();
 using release_t                                         = void (*)(nvid_ctx *);
 using rxHash_t                                          = bool (*)(nvid_ctx *, uint32_t, uint64_t, uint32_t *, uint32_t *);
 using rxPrepare_t                                       = bool (*)(nvid_ctx *, const void *, size_t, bool, uint32_t);
+using rxUpdateDataset_t                                 = bool (*)(nvid_ctx *, const void *, size_t);
 using setJob_t                                          = bool (*)(nvid_ctx *, const void *, size_t, uint32_t);
 using setJob_v2_t                                       = bool (*)(nvid_ctx *, const void *, size_t, const char *);
 using version_t                                         = uint32_t (*)(Version);
@@ -116,6 +118,7 @@ static pluginVersion_t pPluginVersion                   = nullptr;
 static release_t pRelease                               = nullptr;
 static rxHash_t pRxHash                                 = nullptr;
 static rxPrepare_t pRxPrepare                           = nullptr;
+static rxUpdateDataset_t pRxUpdateDataset               = nullptr;
 static setJob_t pSetJob                                 = nullptr;
 static setJob_v2_t pSetJob_v2                           = nullptr;
 static version_t pVersion                               = nullptr;
@@ -202,10 +205,26 @@ bool xmrig::CudaLib::rxHash(nvid_ctx *ctx, uint32_t startNonce, uint64_t target,
 
 bool xmrig::CudaLib::rxPrepare(nvid_ctx *ctx, const void *dataset, size_t datasetSize, bool dataset_host, uint32_t batchSize) noexcept
 {
+#   ifdef XMRIG_ALGO_RANDOMX
+    if (!pRxUpdateDataset) {
+        LOG_WARN("%s" YELLOW_BOLD("CUDA plugin is outdated. Please update to the latest version"), Tags::randomx());
+    }
+#   endif
+
     return pRxPrepare(ctx, dataset, datasetSize, dataset_host, batchSize);
 }
 
 
+bool xmrig::CudaLib::rxUpdateDataset(nvid_ctx *ctx, const void *dataset, size_t datasetSize) noexcept
+{
+    if (pRxUpdateDataset) {
+        return pRxUpdateDataset(ctx, dataset, datasetSize);
+    }
+
+    return true;
+}
+
+
 bool xmrig::CudaLib::kawPowHash(nvid_ctx *ctx, uint8_t* job_blob, uint64_t target, uint32_t *rescount, uint32_t *resnonce, uint32_t *skipped_hashes) noexcept
 {
     return pKawPowHash(ctx, job_blob, target, rescount, resnonce, skipped_hashes);
@@ -401,5 +420,7 @@ void xmrig::CudaLib::load()
         DLSYM(SetJob_v2);
     }
 
+    uv_dlsym(&cudaLib, kRxUpdateDataset, reinterpret_cast<void**>(&pRxUpdateDataset));
+
     pInit();
 }

src/backend/cuda/wrappers/CudaLib.h

@@ -1,6 +1,6 @@
 /* XMRig
- * Copyright (c) 2018-2021 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2021 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright (c) 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright (c) 2016-2024 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
@@ -71,6 +71,7 @@ public:
     static bool deviceInit(nvid_ctx *ctx) noexcept;
     static bool rxHash(nvid_ctx *ctx, uint32_t startNonce, uint64_t target, uint32_t *rescount, uint32_t *resnonce) noexcept;
     static bool rxPrepare(nvid_ctx *ctx, const void *dataset, size_t datasetSize, bool dataset_host, uint32_t batchSize) noexcept;
+    static bool rxUpdateDataset(nvid_ctx *ctx, const void *dataset, size_t datasetSize) noexcept;
     static bool kawPowHash(nvid_ctx *ctx, uint8_t* job_blob, uint64_t target, uint32_t *rescount, uint32_t *resnonce, uint32_t *skipped_hashes) noexcept;
     static bool kawPowPrepare(nvid_ctx *ctx, const void* cache, size_t cache_size, const void* dag_precalc, size_t dag_size, uint32_t height, const uint64_t* dag_sizes) noexcept;
     static bool kawPowStopHash(nvid_ctx *ctx) noexcept;

src/backend/opencl/OclBackend.cpp

@@ -283,7 +283,7 @@ const char *xmrig::ocl_tag()
 
 
 xmrig::OclBackend::OclBackend(Controller *controller) :
-    d_ptr(new OclBackendPrivate(controller))
+    d_ptr(std::make_shared<OclBackendPrivate>(controller))
 {
     d_ptr->workers.setBackend(this);
 }
@@ -291,7 +291,7 @@ xmrig::OclBackend::OclBackend(Controller *controller) :
 
 xmrig::OclBackend::~OclBackend()
 {
-    delete d_ptr;
+    d_ptr.reset();
 
     OclLib::close();
 
@@ -352,15 +352,20 @@ void xmrig::OclBackend::printHashrate(bool details)
 
     char num[16 * 3] = { 0 };
 
-    const double hashrate_short  = hashrate()->calc(Hashrate::ShortInterval);
-    const double hashrate_medium = hashrate()->calc(Hashrate::MediumInterval);
-    const double hashrate_large  = hashrate()->calc(Hashrate::LargeInterval);
+    auto hashrate_short = hashrate()->calc(Hashrate::ShortInterval);
+    auto hashrate_medium = hashrate()->calc(Hashrate::MediumInterval);
+    auto hashrate_large = hashrate()->calc(Hashrate::LargeInterval);
 
     double scale = 1.0;
     const char* h = " H/s";
 
-    if ((hashrate_short >= 1e6) || (hashrate_medium >= 1e6) || (hashrate_large >= 1e6)) {
+    if ((hashrate_short.second >= 1e6) || (hashrate_medium.second >= 1e6) || (hashrate_large.second >= 1e6)) {
         scale = 1e-6;
+
+        hashrate_short.second  *= scale;
+        hashrate_medium.second *= scale;
+        hashrate_large.second  *= scale;
+
         h = "MH/s";
     }
 
@@ -368,12 +373,16 @@ void xmrig::OclBackend::printHashrate(bool details)
 
     size_t i = 0;
     for (const auto& data : d_ptr->threads) {
+        auto h0 = hashrate()->calc(i, Hashrate::ShortInterval);
+        auto h1 = hashrate()->calc(i, Hashrate::MediumInterval);
+        auto h2 = hashrate()->calc(i, Hashrate::LargeInterval);
+
         Log::print("| %8zu | %8" PRId64 " | %8s | %8s | %8s |" CYAN_BOLD(" #%u") YELLOW(" %s") " %s",
                     i,
                     data.affinity,
-                    Hashrate::format(hashrate()->calc(i, Hashrate::ShortInterval)  * scale, num,          sizeof num / 3),
-                    Hashrate::format(hashrate()->calc(i, Hashrate::MediumInterval) * scale, num + 16,     sizeof num / 3),
-                    Hashrate::format(hashrate()->calc(i, Hashrate::LargeInterval)  * scale, num + 16 * 2, sizeof num / 3),
+                    Hashrate::format(h0, num, sizeof num / 3),
+                    Hashrate::format(h1, num + 16, sizeof num / 3),
+                    Hashrate::format(h2, num + 16 * 2, sizeof num / 3),
                     data.device.index(),
                     data.device.topology().toString().data(),
                     data.device.printableName().data()
@@ -383,9 +392,9 @@ void xmrig::OclBackend::printHashrate(bool details)
     }
 
     Log::print(WHITE_BOLD_S "|        - |        - | %8s | %8s | %8s |",
-               Hashrate::format(hashrate_short  * scale, num,          sizeof num / 3),
-               Hashrate::format(hashrate_medium * scale, num + 16,     sizeof num / 3),
-               Hashrate::format(hashrate_large  * scale, num + 16 * 2, sizeof num / 3)
+               Hashrate::format(hashrate_short , num,          sizeof num / 3),
+               Hashrate::format(hashrate_medium, num + 16,     sizeof num / 3),
+               Hashrate::format(hashrate_large , num + 16 * 2, sizeof num / 3)
                );
 }
 

src/backend/opencl/OclBackend.h

@@ -70,7 +70,7 @@ protected:
 #   endif
 
 private:
-    OclBackendPrivate *d_ptr;
+    std::shared_ptr<OclBackendPrivate> d_ptr;
 };
 
 

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()), results);
+                m_runner->run(readUnaligned(m_job.nonce()), m_job.nonceOffset(), results);
             }
             catch (std::exception &ex) {
                 printError(id(), ex.what());
@@ -190,7 +190,7 @@ void xmrig::OclWorker::start()
             std::this_thread::yield();
         }
 
-        if (!consumeJob()) {
+        if (isReady() && !consumeJob()) {
             return;
         }
     }

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

@@ -22,8 +22,8 @@
 #define ALGO_RX_WOW         0x72141177
 #define ALGO_RX_ARQMA       0x72121061
 #define ALGO_RX_SFX         0x72151273
-#define ALGO_RX_KEVA        0x7214116b
 #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

@@ -34,9 +34,9 @@ static const char cryptonight_cl[61447] = {
     0x31,0x35,0x31,0x32,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x57,0x4f,0x57,0x20,0x30,0x78,0x37,0x32,0x31,
     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,0x4b,0x45,0x56,0x41,0x20,0x30,0x78,0x37,0x32,0x31,
-    0x34,0x31,0x31,0x36,0x62,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,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,0x41,0x20,0x30,
+    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,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,6 +225,110 @@ __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,13 +1,11 @@
 #include "../cn/algorithm.cl"
 
-#if (ALGO == ALGO_RX_0)
+#if ((ALGO == ALGO_RX_0) || (ALGO == ALGO_RX_YADA))
 #include "randomx_constants_monero.h"
 #elif (ALGO == ALGO_RX_WOW)
 #include "randomx_constants_wow.h"
 #elif (ALGO == ALGO_RX_ARQMA)
 #include "randomx_constants_arqma.h"
-#elif (ALGO == ALGO_RX_KEVA)
-#include "randomx_constants_keva.h"
 #elif (ALGO == ALGO_RX_GRAFT)
 #include "randomx_constants_graft.h"
 #endif

src/backend/opencl/cl/rx/randomx_constants_keva.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      1048576
-
-//Scratchpad L2 size in bytes. Must be a power of two and less than or equal to RANDOMX_SCRATCHPAD_L3.
-#define RANDOMX_SCRATCHPAD_L2      131072
-
-//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 - 17)
-#define LOC_L3 (32 - 20)

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 *hashOutput)  = 0;
+    virtual void run(uint32_t nonce, uint32_t nonce_offset, 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

@@ -0,0 +1,59 @@
+/* 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

@@ -0,0 +1,50 @@
+/* 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

@@ -5,13 +5,7 @@ if (BUILD_STATIC AND XMRIG_OS_UNIX AND WITH_OPENCL)
 endif()
 
 if (WITH_OPENCL)
-    add_definitions(/DXMRIG_FEATURE_OPENCL)
-    add_definitions(/DCL_USE_DEPRECATED_OPENCL_1_2_APIS)
-    if (XMRIG_OS_APPLE)
-        add_definitions(/DCL_TARGET_OPENCL_VERSION=120)
-    elseif (WITH_OPENCL_VERSION)
-        add_definitions(/DCL_TARGET_OPENCL_VERSION=${WITH_OPENCL_VERSION})
-    endif()
+    add_definitions(/DXMRIG_FEATURE_OPENCL /DCL_USE_DEPRECATED_OPENCL_1_2_APIS)
 
     set(HEADERS_BACKEND_OPENCL
         src/backend/opencl/cl/OclSource.h
@@ -71,6 +65,13 @@ if (WITH_OPENCL)
         src/backend/opencl/wrappers/OclPlatform.cpp
         )
 
+    if (XMRIG_OS_APPLE)
+        add_definitions(/DCL_TARGET_OPENCL_VERSION=120)
+        list(APPEND SOURCES_BACKEND_OPENCL src/backend/opencl/wrappers/OclDevice_mac.cpp)
+    elseif (WITH_OPENCL_VERSION)
+        add_definitions(/DCL_TARGET_OPENCL_VERSION=${WITH_OPENCL_VERSION})
+    endif()
+
     if (WIN32)
         list(APPEND SOURCES_BACKEND_OPENCL src/backend/opencl/OclCache_win.cpp)
     else()
@@ -80,6 +81,7 @@ 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
@@ -97,6 +99,7 @@ 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 *hashOutput)
+void xmrig::OclCnRunner::run(uint32_t nonce, uint32_t /*nonce_offset*/, 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 *hashOutput) override;
+    void run(uint32_t nonce, uint32_t nonce_offset, 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 *hashOutput)
+void OclKawPowRunner::run(uint32_t nonce, uint32_t /*nonce_offset*/, 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 *hashOutput) override;
+    void run(uint32_t nonce, uint32_t nonce_offset, 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,6 +26,7 @@
 #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"
@@ -73,6 +74,7 @@ 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;
@@ -85,7 +87,7 @@ xmrig::OclRxBaseRunner::~OclRxBaseRunner()
 }
 
 
-void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
+void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput)
 {
     static const uint32_t zero = 0;
 
@@ -96,8 +98,7 @@ void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
         m_blake2b_initial_hash_double->setNonce(nonce);
     }
     else {
-        hashOutput[0xFF] = 0;
-        return;
+        m_blake2b_initial_hash_big->setNonce(nonce, nonce_offset);
     }
 
     m_find_shares->setNonce(nonce);
@@ -107,9 +108,12 @@ void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
     if (m_jobSize <= 128) {
         m_blake2b_initial_hash->enqueue(m_queue, m_intensity);
     }
-    else {
+    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_fillAes1Rx4_scratchpad->enqueue(m_queue, m_intensity);
 
@@ -150,12 +154,15 @@ 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());
 }
@@ -191,6 +198,9 @@ 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,6 +36,7 @@ namespace xmrig {
 class Blake2bHashRegistersKernel;
 class Blake2bInitialHashKernel;
 class Blake2bInitialHashDoubleKernel;
+class Blake2bInitialHashBigKernel;
 class FillAesKernel;
 class FindSharesKernel;
 class HashAesKernel;
@@ -53,7 +54,7 @@ protected:
     size_t bufferSize() const override;
     void build() override;
     void init() override;
-    void run(uint32_t nonce, uint32_t *hashOutput) override;
+    void run(uint32_t nonce, uint32_t nonce_offset, uint32_t *hashOutput) override;
     void set(const Job &job, uint8_t *blob) override;
 
 protected:
@@ -63,6 +64,7 @@ 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/backend/opencl/wrappers/OclDevice.cpp

@@ -1,6 +1,7 @@
 /* XMRig
- * Copyright (c) 2018-2021 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2021 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright (c) 2021      Spudz76     <https://github.com/Spudz76>
+ * Copyright (c) 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright (c) 2016-2024 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
@@ -17,6 +18,7 @@
  */
 
 #include "backend/opencl/wrappers/OclDevice.h"
+#include "3rdparty/fmt/core.h"
 #include "3rdparty/rapidjson/document.h"
 #include "backend/opencl/OclGenerator.h"
 #include "backend/opencl/OclThreads.h"
@@ -30,19 +32,21 @@
 
 
 #include <algorithm>
-
-
-// NOLINTNEXTLINE(modernize-use-using)
-typedef union
-{
-    struct { cl_uint type; cl_uint data[5]; } raw;
-    struct { cl_uint type; cl_char unused[17]; cl_char bus; cl_char device; cl_char function; } pcie;
-} topology_amd;
+#include <map>
 
 
 namespace xmrig {
 
 
+struct topology_amd {
+    cl_uint type;
+    cl_char unused[17];
+    cl_char bus;
+    cl_char device;
+    cl_char function;
+};
+
+
 #ifdef XMRIG_ALGO_RANDOMX
 extern bool ocl_generic_rx_generator(const OclDevice &device, const Algorithm &algorithm, OclThreads &threads);
 #endif
@@ -81,9 +85,11 @@ static OclVendor getPlatformVendorId(const String &vendor, const String &extensi
         return OCL_VENDOR_INTEL;
     }
 
+#   ifdef XMRIG_OS_APPLE
     if (extensions.contains("cl_APPLE_") || vendor.contains("Apple")) {
         return OCL_VENDOR_APPLE;
     }
+#   endif
 
     return OCL_VENDOR_UNKNOWN;
 }
@@ -103,117 +109,16 @@ static OclVendor getVendorId(const String &vendor)
         return OCL_VENDOR_INTEL;
     }
 
+#   ifdef XMRIG_OS_APPLE
     if (vendor.contains("Apple")) {
         return OCL_VENDOR_APPLE;
     }
+#   endif
 
     return OCL_VENDOR_UNKNOWN;
 }
 
 
-static OclDevice::Type getType(const String &name, const OclVendor platformVendorId)
-{
-    if (platformVendorId == OCL_VENDOR_APPLE) {
-        // Apple Platform: uses product names, not gfx# or codenames
-        if (name.contains("AMD Radeon")) {
-            if (name.contains(" 450 ") ||
-                name.contains(" 455 ") ||
-                name.contains(" 460 ")) {
-                return OclDevice::Baffin;
-            }
-
-            if (name.contains(" 555 ") || name.contains(" 555X ") ||
-                name.contains(" 560 ") || name.contains(" 560X ") ||
-                name.contains(" 570 ") || name.contains(" 570X ") ||
-                name.contains(" 575 ") || name.contains(" 575X ")) {
-                return OclDevice::Polaris;
-            }
-
-            if (name.contains(" 580 ") || name.contains(" 580X ")) {
-                return OclDevice::Ellesmere;
-            }
-
-            if (name.contains(" Vega ")) {
-                if (name.contains(" 48 ") ||
-                    name.contains(" 56 ") ||
-                    name.contains(" 64 ") ||
-                    name.contains(" 64X ")) {
-                    return OclDevice::Vega_10;
-                }
-                if (name.contains(" 16 ") ||
-                    name.contains(" 20 ") ||
-                    name.contains(" II ")) {
-                    return OclDevice::Vega_20;
-                }
-            }
-
-            if (name.contains(" 5700 ") || name.contains(" W5700X ")) {
-                return OclDevice::Navi_10;
-            }
-
-            if (name.contains(" 5600 ") || name.contains(" 5600M ")) {
-                return OclDevice::Navi_12;
-            }
-
-            if (name.contains(" 5300 ") || name.contains(" 5300M ") ||
-                name.contains(" 5500 ") || name.contains(" 5500M ")) {
-                return OclDevice::Navi_14;
-            }
-
-            if (name.contains(" W6800 ") || name.contains(" W6900X ")) {
-                return OclDevice::Navi_21;
-            }
-        }
-    }
-
-    if (name == "gfx900" || name == "gfx901") {
-        return OclDevice::Vega_10;
-    }
-
-    if (name == "gfx902" || name == "gfx903") {
-        return OclDevice::Raven;
-    }
-
-    if (name == "gfx906" || name == "gfx907") {
-        return OclDevice::Vega_20;
-    }
-
-    if (name == "gfx1010") {
-        return OclDevice::Navi_10;
-    }
-
-    if (name == "gfx1011") {
-        return OclDevice::Navi_12;
-    }
-
-    if (name == "gfx1012") {
-        return OclDevice::Navi_14;
-    }
-
-    if (name == "gfx1030") {
-        return OclDevice::Navi_21;
-    }
-
-    if (name == "gfx804") {
-        return OclDevice::Lexa;
-    }
-
-    if (name == "Baffin") {
-        return OclDevice::Baffin;
-    }
-
-    if (name.contains("Ellesmere")) {
-        return OclDevice::Ellesmere;
-    }
-
-    if (name == "gfx803" || name.contains("polaris")) {
-        return OclDevice::Polaris;
-    }
-
-    return OclDevice::Unknown;
-}
-
-
 } // namespace xmrig
 
 
@@ -231,21 +136,21 @@ xmrig::OclDevice::OclDevice(uint32_t index, cl_device_id id, cl_platform_id plat
 {
     m_vendorId  = getVendorId(m_vendor);
     m_platformVendorId = getPlatformVendorId(m_platformVendor, m_extensions);
-    m_type      = getType(m_name, m_platformVendorId);
+    m_type      = getType(m_name);
 
     if (m_extensions.contains("cl_amd_device_attribute_query")) {
-        topology_amd topology;
-
-        if (OclLib::getDeviceInfo(id, CL_DEVICE_TOPOLOGY_AMD, sizeof(topology), &topology, nullptr) == CL_SUCCESS && topology.raw.type == CL_DEVICE_TOPOLOGY_TYPE_PCIE_AMD) {
-            m_topology = PciTopology(static_cast<uint32_t>(topology.pcie.bus), static_cast<uint32_t>(topology.pcie.device), static_cast<uint32_t>(topology.pcie.function));
+        topology_amd topology{};
+        if (OclLib::getDeviceInfo(id, CL_DEVICE_TOPOLOGY_AMD, sizeof(topology), &topology) == CL_SUCCESS && topology.type == CL_DEVICE_TOPOLOGY_TYPE_PCIE_AMD) {
+            m_topology = { topology.bus, topology.device, topology.function };
         }
+
         m_board = OclLib::getString(id, CL_DEVICE_BOARD_NAME_AMD);
     }
     else if (m_extensions.contains("cl_nv_device_attribute_query")) {
         cl_uint bus = 0;
-        if (OclLib::getDeviceInfo(id, CL_DEVICE_PCI_BUS_ID_NV, sizeof (bus), &bus, nullptr) == CL_SUCCESS) {
+        if (OclLib::getDeviceInfo(id, CL_DEVICE_PCI_BUS_ID_NV, sizeof(bus), &bus) == CL_SUCCESS) {
             cl_uint slot  = OclLib::getUint(id, CL_DEVICE_PCI_SLOT_ID_NV);
-            m_topology = PciTopology(bus, (slot >> 3) & 0xff, slot & 7);
+            m_topology = { bus, (slot >> 3) & 0xff, slot & 7 };
         }
     }
 }
@@ -253,17 +158,11 @@ xmrig::OclDevice::OclDevice(uint32_t index, cl_device_id id, cl_platform_id plat
 
 xmrig::String xmrig::OclDevice::printableName() const
 {
-    const size_t size = m_board.size() + m_name.size() + 64;
-    char *buf         = new char[size]();
-
     if (m_board.isNull()) {
-        snprintf(buf, size, GREEN_BOLD("%s"), m_name.data());
-    }
-    else {
-        snprintf(buf, size, GREEN_BOLD("%s") " (" CYAN_BOLD("%s") ")", m_board.data(), m_name.data());
+        return fmt::format(GREEN_BOLD("{}"), m_name).c_str();
     }
 
-    return buf;
+    return fmt::format(GREEN_BOLD("{}") " (" CYAN_BOLD("{}") ")", m_board, m_name).c_str();
 }
 
 
@@ -311,3 +210,35 @@ void xmrig::OclDevice::toJSON(rapidjson::Value &out, rapidjson::Document &doc) c
 #   endif
 }
 #endif
+
+
+#ifndef XMRIG_OS_APPLE
+xmrig::OclDevice::Type xmrig::OclDevice::getType(const String &name)
+{
+    static std::map<const char *, OclDevice::Type> types = {
+        { "gfx900",     Vega_10 },
+        { "gfx901",     Vega_10 },
+        { "gfx902",     Raven },
+        { "gfx903",     Raven },
+        { "gfx906",     Vega_20 },
+        { "gfx907",     Vega_20 },
+        { "gfx1010",    Navi_10 },
+        { "gfx1011",    Navi_12 },
+        { "gfx1012",    Navi_14 },
+        { "gfx1030",    Navi_21 },
+        { "gfx804",     Lexa },
+        { "Baffin",     Baffin },
+        { "Ellesmere",  Ellesmere },
+        { "gfx803",     Polaris },
+        { "polaris",    Polaris },
+    };
+
+    for (auto &kv : types) {
+        if (name.contains(kv.first)) {
+            return kv.second;
+        }
+    }
+
+    return OclDevice::Unknown;
+}
+#endif

src/backend/opencl/wrappers/OclDevice.h

@@ -1,6 +1,6 @@
 /* XMRig
- * Copyright (c) 2018-2021 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2021 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright (c) 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright (c) 2016-2024 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
@@ -86,6 +86,8 @@ public:
 #   endif
 
 private:
+    static OclDevice::Type getType(const String &name);
+
     cl_device_id m_id               = nullptr;
     cl_platform_id m_platform       = nullptr;
     const String m_platformVendor;

src/backend/opencl/wrappers/OclDevice_mac.cpp

@@ -0,0 +1,77 @@
+/* XMRig
+ * Copyright (c) 2021      Spudz76     <https://github.com/Spudz76>
+ * Copyright (c) 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright (c) 2016-2024 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ *
+ *   This program is free software: you can redistribute it and/or modify
+ *   it under the terms of the GNU General Public License as published by
+ *   the Free Software Foundation, either version 3 of the License, or
+ *   (at your option) any later version.
+ *
+ *   This program is distributed in the hope that it will be useful,
+ *   but WITHOUT ANY WARRANTY; without even the implied warranty of
+ *   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ *   GNU General Public License for more details.
+ *
+ *   You should have received a copy of the GNU General Public License
+ *   along with this program. If not, see <http://www.gnu.org/licenses/>.
+ */
+
+#include "backend/opencl/wrappers/OclDevice.h"
+
+
+xmrig::OclDevice::Type xmrig::OclDevice::getType(const String &name)
+{
+    // Apple Platform: uses product names, not gfx# or codenames
+    if (name.contains("AMD Radeon")) {
+        if (name.contains(" 450 ") ||
+            name.contains(" 455 ") ||
+            name.contains(" 460 ")) {
+            return Baffin;
+        }
+
+        if (name.contains(" 555 ") || name.contains(" 555X ") ||
+            name.contains(" 560 ") || name.contains(" 560X ") ||
+            name.contains(" 570 ") || name.contains(" 570X ") ||
+            name.contains(" 575 ") || name.contains(" 575X ")) {
+            return Polaris;
+        }
+
+        if (name.contains(" 580 ") || name.contains(" 580X ")) {
+            return Ellesmere;
+        }
+
+        if (name.contains(" Vega ")) {
+            if (name.contains(" 48 ") ||
+                name.contains(" 56 ") ||
+                name.contains(" 64 ") ||
+                name.contains(" 64X ")) {
+                return Vega_10;
+            }
+            if (name.contains(" 16 ") ||
+                name.contains(" 20 ") ||
+                name.contains(" II ")) {
+                return Vega_20;
+            }
+        }
+
+        if (name.contains(" 5700 ") || name.contains(" W5700X ")) {
+            return Navi_10;
+        }
+
+        if (name.contains(" 5600 ") || name.contains(" 5600M ")) {
+            return Navi_12;
+        }
+
+        if (name.contains(" 5300 ") || name.contains(" 5300M ") ||
+            name.contains(" 5500 ") || name.contains(" 5500M ")) {
+            return Navi_14;
+        }
+
+        if (name.contains(" W6800 ") || name.contains(" W6900X ")) {
+            return Navi_21;
+        }
+    }
+
+    return OclDevice::Unknown;
+}

src/base/api/Api.cpp

@@ -95,8 +95,7 @@ xmrig::Api::~Api()
 #   ifdef XMRIG_FEATURE_HTTP
     if (m_httpd) {
         m_httpd->stop();
-        delete m_httpd;
-        m_httpd = nullptr; // Ensure the pointer is set to nullptr after deletion
+        m_httpd.reset();
     }
 #   endif
 }
@@ -116,12 +115,11 @@ void xmrig::Api::start()
 
 #   ifdef XMRIG_FEATURE_HTTP
     if (!m_httpd) {
-        m_httpd = new Httpd(m_base);
+        m_httpd = std::make_shared<Httpd>(m_base);
         if (!m_httpd->start()) {
             LOG_ERR("%s " RED_BOLD("HTTP API server failed to start."), Tags::network());
 
-            delete m_httpd; // Properly handle failure to start
-            m_httpd = nullptr;
+            m_httpd.reset();
         }
     }
 #   endif

src/base/api/Api.h

@@ -66,7 +66,7 @@ private:
     Base *m_base;
     char m_id[32]{};
     const uint64_t m_timestamp;
-    Httpd *m_httpd  = nullptr;
+    std::shared_ptr<Httpd> m_httpd;
     std::vector<IApiListener *> m_listeners;
     String m_workerId;
     uint8_t m_ticks = 0;

src/base/api/Httpd.cpp

@@ -69,13 +69,13 @@ bool xmrig::Httpd::start()
     bool tls = false;
 
 #   ifdef XMRIG_FEATURE_TLS
-    m_http = new HttpsServer(m_httpListener);
+    m_http = std::make_shared<HttpsServer>(m_httpListener);
     tls = m_http->setTls(m_base->config()->tls());
 #   else
-    m_http = new HttpServer(m_httpListener);
+    m_http = std::make_shared<HttpServer>(m_httpListener);
 #   endif
 
-    m_server = new TcpServer(config.host(), config.port(), m_http);
+    m_server = std::make_shared<TcpServer>(config.host(), config.port(), m_http.get());
 
     const int rc = m_server->bind();
     Log::print(GREEN_BOLD(" * ") WHITE_BOLD("%-13s") CSI "1;%dm%s:%d" " " RED_BOLD("%s"),
@@ -112,9 +112,6 @@ bool xmrig::Httpd::start()
 
 void xmrig::Httpd::stop()
 {
-    delete m_server;
-    delete m_http;
-
     m_server = nullptr;
     m_http   = nullptr;
     m_port   = 0;

src/base/api/Httpd.h

@@ -55,13 +55,13 @@ private:
 
     const Base *m_base;
     std::shared_ptr<IHttpListener> m_httpListener;
-    TcpServer *m_server     = nullptr;
+    std::shared_ptr<TcpServer> m_server;
     uint16_t m_port         = 0;
 
 #   ifdef XMRIG_FEATURE_TLS
-    HttpsServer *m_http     = nullptr;
+    std::shared_ptr<HttpsServer> m_http;
 #   else
-    HttpServer *m_http      = nullptr;
+    std::shared_ptr<HttpServer> m_http;
 #   endif
 };
 

src/base/crypto/Algorithm.cpp

@@ -81,7 +81,7 @@ const char *Algorithm::kRX_WOW          = "rx/wow";
 const char *Algorithm::kRX_ARQ          = "rx/arq";
 const char *Algorithm::kRX_GRAFT        = "rx/graft";
 const char *Algorithm::kRX_SFX          = "rx/sfx";
-const char *Algorithm::kRX_KEVA         = "rx/keva";
+const char *Algorithm::kRX_YADA         = "rx/yada";
 #endif
 
 #ifdef XMRIG_ALGO_ARGON2
@@ -147,7 +147,7 @@ static const std::map<uint32_t, const char *> kAlgorithmNames = {
     ALGO_NAME(RX_ARQ),
     ALGO_NAME(RX_GRAFT),
     ALGO_NAME(RX_SFX),
-    ALGO_NAME(RX_KEVA),
+    ALGO_NAME(RX_YADA),
 #   endif
 
 #   ifdef XMRIG_ALGO_ARGON2
@@ -261,8 +261,8 @@ static const std::map<const char *, Algorithm::Id, aliasCompare> kAlgorithmAlias
                                     ALGO_ALIAS(RX_GRAFT,        "randomgraft"),
     ALGO_ALIAS_AUTO(RX_SFX),        ALGO_ALIAS(RX_SFX,          "randomx/sfx"),
                                     ALGO_ALIAS(RX_SFX,          "randomsfx"),
-    ALGO_ALIAS_AUTO(RX_KEVA),       ALGO_ALIAS(RX_KEVA,         "randomx/keva"),
-                                    ALGO_ALIAS(RX_KEVA,         "randomkeva"),
+    ALGO_ALIAS_AUTO(RX_YADA),       ALGO_ALIAS(RX_YADA,         "randomx/yada"),
+                                    ALGO_ALIAS(RX_YADA,         "randomyada"),
 #   endif
 
 #   ifdef XMRIG_ALGO_ARGON2
@@ -350,7 +350,7 @@ std::vector<xmrig::Algorithm> xmrig::Algorithm::all(const std::function<bool(con
         CN_HEAVY_0, CN_HEAVY_TUBE, CN_HEAVY_XHV,
         CN_PICO_0, CN_PICO_TLO,
         CN_UPX2,
-        RX_0, RX_WOW, RX_ARQ, RX_GRAFT, RX_SFX, RX_KEVA,
+        RX_0, RX_WOW, RX_ARQ, RX_GRAFT, RX_SFX, RX_YADA,
         AR2_CHUKWA, AR2_CHUKWA_V2, AR2_WRKZ,
         KAWPOW_RVN,
         GHOSTRIDER_RTM

src/base/crypto/Algorithm.h

@@ -77,7 +77,7 @@ public:
         RX_ARQ          = 0x72121061,   // "rx/arq"           RandomARQ (Arqma).
         RX_GRAFT        = 0x72151267,   // "rx/graft"         RandomGRAFT (Graft).
         RX_SFX          = 0x72151273,   // "rx/sfx"           RandomSFX (Safex Cash).
-        RX_KEVA         = 0x7214116b,   // "rx/keva"          RandomKEVA (Keva).
+        RX_YADA         = 0x72151279,   // "rx/yada"          RandomYada (YadaCoin).
         AR2_CHUKWA      = 0x61130000,   // "argon2/chukwa"    Argon2id (Chukwa).
         AR2_CHUKWA_V2   = 0x61140000,   // "argon2/chukwav2"  Argon2id (Chukwa v2).
         AR2_WRKZ        = 0x61120000,   // "argon2/wrkz"      Argon2id (WRKZ)
@@ -143,7 +143,7 @@ public:
     static const char *kRX_ARQ;
     static const char *kRX_GRAFT;
     static const char *kRX_SFX;
-    static const char *kRX_KEVA;
+    static const char *kRX_YADA;
 #   endif
 
 #   ifdef XMRIG_ALGO_ARGON2

src/base/crypto/Coin.cpp

@@ -50,11 +50,11 @@ static const CoinInfo coinInfo[] = {
     { Algorithm::CN_R,            "SUMO",     "Sumokoin",     240,    1000000000,     BLUE_BG_BOLD(   WHITE_BOLD_S " sumo    ") },
     { Algorithm::RX_ARQ,          "ARQ",      "ArQmA",        120,    1000000000,     BLUE_BG_BOLD(   WHITE_BOLD_S " arqma   ") },
     { Algorithm::RX_GRAFT,        "GRFT",     "Graft",        120,    10000000000,    BLUE_BG_BOLD(   WHITE_BOLD_S " graft   ") },
-    { Algorithm::RX_KEVA,         "KVA",      "Kevacoin",     0,      0,              MAGENTA_BG_BOLD(WHITE_BOLD_S " keva    ") },
     { Algorithm::KAWPOW_RVN,      "RVN",      "Ravencoin",    0,      0,              BLUE_BG_BOLD(   WHITE_BOLD_S " raven   ") },
     { Algorithm::RX_WOW,          "WOW",      "Wownero",      300,    100000000000,   MAGENTA_BG_BOLD(WHITE_BOLD_S " wownero ") },
     { Algorithm::RX_0,            "ZEPH",     "Zephyr",       120,    1000000000000,  BLUE_BG_BOLD(   WHITE_BOLD_S " zephyr  ") },
     { Algorithm::RX_0,            "Townforge","Townforge",    30,     100000000,      MAGENTA_BG_BOLD(WHITE_BOLD_S " townforge ") },
+    { Algorithm::RX_YADA,         "YDA",      "YadaCoin",     120,    100000000,      BLUE_BG_BOLD(   WHITE_BOLD_S " yada    ") },
 };
 
 

src/base/crypto/Coin.h

@@ -36,11 +36,11 @@ public:
         SUMO,
         ARQMA,
         GRAFT,
-        KEVA,
         RAVEN,
         WOWNERO,
         ZEPHYR,
         TOWNFORGE,
+        YADA,
         MAX
     };
 

src/base/io/Async.cpp

@@ -128,7 +128,7 @@ public:
 } // namespace xmrig
 
 
-xmrig::Async::Async(Callback callback) : d_ptr(new AsyncPrivate())
+xmrig::Async::Async(Callback callback) : d_ptr(std::make_shared<AsyncPrivate>())
 {
     d_ptr->callback     = std::move(callback);
     d_ptr->async        = new uv_async_t;
@@ -151,8 +151,6 @@ xmrig::Async::Async(IAsyncListener *listener) : d_ptr(new AsyncPrivate())
 xmrig::Async::~Async()
 {
     Handle::close(d_ptr->async);
-
-    delete d_ptr;
 }
 
 

src/base/io/Async.h

@@ -49,7 +49,7 @@ public:
     void send();
 
 private:
-    AsyncPrivate *d_ptr;
+    std::shared_ptr<AsyncPrivate> d_ptr;
 };
 
 

src/base/io/Watcher.cpp

@@ -36,7 +36,7 @@ xmrig::Watcher::Watcher(const String &path, IWatcherListener *listener) :
     m_listener(listener),
     m_path(path)
 {
-    m_timer = new Timer(this);
+    m_timer = std::make_shared<Timer>(this);
 
     m_fsEvent = new uv_fs_event_t;
     m_fsEvent->data = this;
@@ -48,8 +48,6 @@ xmrig::Watcher::Watcher(const String &path, IWatcherListener *listener) :
 
 xmrig::Watcher::~Watcher()
 {
-    delete m_timer;
-
     Handle::close(m_fsEvent);
 }
 

src/base/io/Watcher.h

@@ -60,7 +60,7 @@ private:
 
     IWatcherListener *m_listener;
     String m_path;
-    Timer *m_timer;
+    std::shared_ptr<Timer> m_timer;
     uv_fs_event_t *m_fsEvent;
 };
 

src/base/io/json/Json.cpp

@@ -211,11 +211,13 @@ rapidjson::Value xmrig::Json::normalize(double value, bool zero)
 {
     using namespace rapidjson;
 
-    if (!std::isnormal(value)) {
+    const double value_rounded = floor(value * 100.0) / 100.0;
+
+    if (!std::isnormal(value) || !std::isnormal(value_rounded)) {
         return zero ? Value(0.0) : Value(kNullType);
     }
 
-    return Value(floor(value * 100.0) / 100.0);
+    return Value(value_rounded);
 }
 
 

src/base/io/log/Log.cpp

@@ -66,17 +66,10 @@ public:
 
 
     LogPrivate() = default;
+    ~LogPrivate() = default;
 
 
-    inline ~LogPrivate()
-    {
-        for (auto backend : m_backends) {
-            delete backend;
-        }
-    }
-
-
-    inline void add(ILogBackend *backend) { m_backends.push_back(backend); }
+    inline void add(std::shared_ptr<ILogBackend> backend) { m_backends.emplace_back(backend); }
 
 
     void print(Log::Level level, const char *fmt, va_list args)
@@ -108,7 +101,7 @@ public:
         }
 
         if (!m_backends.empty()) {
-            for (auto backend : m_backends) {
+            for (auto& backend : m_backends) {
                 backend->print(ts, level, m_buf, offset, size, true);
                 backend->print(ts, level, txt.c_str(), offset ? (offset - 11) : 0, txt.size(), false);
             }
@@ -188,13 +181,13 @@ private:
 
     char m_buf[Log::kMaxBufferSize]{};
     std::mutex m_mutex;
-    std::vector<ILogBackend*> m_backends;
+    std::vector<std::shared_ptr<ILogBackend>> m_backends;
 };
 
 
 bool Log::m_background      = false;
 bool Log::m_colors          = true;
-LogPrivate *Log::d          = nullptr;
+std::shared_ptr<LogPrivate> Log::d{};
 uint32_t Log::m_verbose     = 0;
 
 
@@ -202,7 +195,7 @@ uint32_t Log::m_verbose     = 0;
 
 
 
-void xmrig::Log::add(ILogBackend *backend)
+void xmrig::Log::add(std::shared_ptr<ILogBackend> backend)
 {
     assert(d != nullptr);
 
@@ -214,14 +207,13 @@ void xmrig::Log::add(ILogBackend *backend)
 
 void xmrig::Log::destroy()
 {
-    delete d;
-    d = nullptr;
+    d.reset();
 }
 
 
 void xmrig::Log::init()
 {
-    d = new LogPrivate();
+    d = std::make_shared<LogPrivate>();
 }
 
 

src/base/io/log/Log.h

@@ -23,6 +23,7 @@
 
 #include <cstddef>
 #include <cstdint>
+#include <memory>
 
 
 namespace xmrig {
@@ -49,7 +50,7 @@ public:
 
     constexpr static size_t kMaxBufferSize = 16384;
 
-    static void add(ILogBackend *backend);
+    static void add(std::shared_ptr<ILogBackend> backend);
     static void destroy();
     static void init();
     static void print(const char *fmt, ...);
@@ -66,7 +67,7 @@ public:
 private:
     static bool m_background;
     static bool m_colors;
-    static LogPrivate *d;
+    static std::shared_ptr<LogPrivate> d;
     static uint32_t m_verbose;
 };
 

src/base/kernel/Base.cpp

@@ -80,11 +80,10 @@ public:
     inline ~BasePrivate()
     {
 #       ifdef XMRIG_FEATURE_API
-        delete api;
+        api.reset();
 #       endif
 
-        delete config;
-        delete watcher;
+        watcher.reset();
 
         NetBuffer::destroy();
     }
@@ -98,27 +97,25 @@ public:
     }
 
 
-    inline void replace(Config *newConfig)
+    inline void replace(std::shared_ptr<Config> newConfig)
     {
-        Config *previousConfig = config;
+        auto previousConfig = config;
         config = newConfig;
 
         for (IBaseListener *listener : listeners) {
-            listener->onConfigChanged(config, previousConfig);
+            listener->onConfigChanged(config.get(), previousConfig.get());
         }
-
-        delete previousConfig;
     }
 
 
-    Api *api            = nullptr;
-    Config *config      = nullptr;
+    std::shared_ptr<Api> api;
+    std::shared_ptr<Config> config;
     std::vector<IBaseListener *> listeners;
-    Watcher *watcher    = nullptr;
+    std::shared_ptr<Watcher> watcher;
 
 
 private:
-    inline static Config *load(Process *process)
+    inline static std::shared_ptr<Config> load(Process *process)
     {
         JsonChain chain;
         ConfigTransform transform;
@@ -127,29 +124,29 @@ private:
         ConfigTransform::load(chain, process, transform);
 
         if (read(chain, config)) {
-            return config.release();
+            return config;
         }
 
         chain.addFile(Process::location(Process::DataLocation, "config.json"));
         if (read(chain, config)) {
-            return config.release();
+            return config;
         }
 
         chain.addFile(Process::location(Process::HomeLocation,  "." APP_ID ".json"));
         if (read(chain, config)) {
-            return config.release();
+            return config;
         }
 
         chain.addFile(Process::location(Process::HomeLocation, ".config" XMRIG_DIR_SEPARATOR APP_ID ".json"));
         if (read(chain, config)) {
-            return config.release();
+            return config;
         }
 
 #       ifdef XMRIG_FEATURE_EMBEDDED_CONFIG
         chain.addRaw(default_config);
 
         if (read(chain, config)) {
-            return config.release();
+            return config;
         }
 #       endif
 
@@ -162,7 +159,7 @@ private:
 
 
 xmrig::Base::Base(Process *process)
-    : d_ptr(new BasePrivate(process))
+    : d_ptr(std::make_shared<BasePrivate>(process))
 {
 
 }
@@ -170,7 +167,6 @@ xmrig::Base::Base(Process *process)
 
 xmrig::Base::~Base()
 {
-    delete d_ptr;
 }
 
 
@@ -183,7 +179,7 @@ bool xmrig::Base::isReady() const
 int xmrig::Base::init()
 {
 #   ifdef XMRIG_FEATURE_API
-    d_ptr->api = new Api(this);
+    d_ptr->api = std::make_shared<Api>(this);
     d_ptr->api->addListener(this);
 #   endif
 
@@ -193,16 +189,16 @@ int xmrig::Base::init()
         Log::setBackground(true);
     }
     else {
-        Log::add(new ConsoleLog(config()->title()));
+        Log::add(std::make_shared<ConsoleLog>(config()->title()));
     }
 
     if (config()->logFile()) {
-        Log::add(new FileLog(config()->logFile()));
+        Log::add(std::make_shared<FileLog>(config()->logFile()));
     }
 
 #   ifdef HAVE_SYSLOG_H
     if (config()->isSyslog()) {
-        Log::add(new SysLog());
+        Log::add(std::make_shared<SysLog>());
     }
 #   endif
 
@@ -221,7 +217,7 @@ void xmrig::Base::start()
     }
 
     if (config()->isWatch()) {
-        d_ptr->watcher = new Watcher(config()->fileName(), this);
+        d_ptr->watcher = std::make_shared<Watcher>(config()->fileName(), this);
     }
 }
 
@@ -232,8 +228,7 @@ void xmrig::Base::stop()
     api()->stop();
 #   endif
 
-    delete d_ptr->watcher;
-    d_ptr->watcher = nullptr;
+    d_ptr->watcher.reset();
 }
 
 
@@ -241,7 +236,7 @@ xmrig::Api *xmrig::Base::api() const
 {
     assert(d_ptr->api != nullptr);
 
-    return d_ptr->api;
+    return d_ptr->api.get();
 }
 
 
@@ -258,18 +253,14 @@ bool xmrig::Base::reload(const rapidjson::Value &json)
         return false;
     }
 
-    auto config = new Config();
+    auto config = std::make_shared<Config>();
     if (!config->read(reader, d_ptr->config->fileName())) {
-        delete config;
-
         return false;
     }
 
     const bool saved = config->save();
 
     if (config->isWatch() && d_ptr->watcher && saved) {
-        delete config;
-
         return true;
     }
 
@@ -281,9 +272,9 @@ bool xmrig::Base::reload(const rapidjson::Value &json)
 
 xmrig::Config* xmrig::Base::config() const
 {
-    assert(d_ptr->config != nullptr);
+    assert(d_ptr->config);
 
-    return d_ptr->config;
+    return d_ptr->config.get();
 }
 
 
@@ -300,12 +291,10 @@ void xmrig::Base::onFileChanged(const String &fileName)
     JsonChain chain;
     chain.addFile(fileName);
 
-    auto config = new Config();
+    auto config = std::make_shared<Config>();
 
     if (!config->read(chain, chain.fileName())) {
         LOG_ERR("%s " RED("reloading failed"), Tags::config());
-
-        delete config;
         return;
     }
 

src/base/kernel/Base.h

@@ -64,7 +64,7 @@ protected:
 #   endif
 
 private:
-    BasePrivate *d_ptr;
+    std::shared_ptr<BasePrivate> d_ptr;
 };
 
 

src/base/kernel/Entry.cpp

@@ -5,8 +5,8 @@
  * 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>
+ * Copyright 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2024 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
@@ -22,11 +22,9 @@
  *   along with this program. If not, see <http://www.gnu.org/licenses/>.
  */
 
-
 #include <cstdio>
 #include <uv.h>
 
-
 #ifdef XMRIG_FEATURE_TLS
 #   include <openssl/opensslv.h>
 #endif
@@ -66,13 +64,13 @@ static int showVersion()
 #   endif
 
     printf("\n features:"
-#   if defined(__i386__) || defined(_M_IX86)
-    " 32-bit"
-#   elif defined(__x86_64__) || defined(_M_AMD64)
+#   if defined(__x86_64__) || defined(_M_AMD64) || defined (__arm64__) || defined (__aarch64__)
     " 64-bit"
+#   else
+    " 32-bit"
 #   endif
 
-#   if defined(__AES__) || defined(_MSC_VER)
+#   if defined(__AES__) || defined(_MSC_VER) || defined(__ARM_FEATURE_CRYPTO)
     " AES"
 #   endif
     "\n");

src/base/net/dns/DnsUvBackend.cpp

@@ -29,13 +29,13 @@
 namespace xmrig {
 
 
-static Storage<DnsUvBackend> *storage = nullptr;
+static std::shared_ptr<Storage<DnsUvBackend>> storage = nullptr;
 
 
 Storage<DnsUvBackend> &DnsUvBackend::getStorage()
 {
-    if (storage == nullptr) {
-        storage = new Storage<DnsUvBackend>();
+    if (!storage) {
+        storage = std::make_shared<Storage<DnsUvBackend>>();
     }
 
     return *storage;
@@ -67,8 +67,7 @@ xmrig::DnsUvBackend::~DnsUvBackend()
     storage->release(m_key);
 
     if (storage->isEmpty()) {
-        delete storage;
-        storage = nullptr;
+        storage.reset();
     }
 }