v4.2.1-beta

Workaround for a bug in binutils-2.32-1 on ARM

CHANGELOG.md

@@ -1,8 +1,42 @@
+# v4.2.1-beta
+- [#1150](https://github.com/xmrig/xmrig/issues/1150) Fixed build on FreeBSD.
+- [#1175](https://github.com/xmrig/xmrig/issues/1175) Fixed support for systems where total count of NUMA nodes not equal usable nodes count.
+- [#1199](https://github.com/xmrig/xmrig/issues/1199) Fixed excessive memory allocation for OpenCL threads with low intensity.
+- [#1212](https://github.com/xmrig/xmrig/issues/1212) Fixed low RandomX performance after fast algorithm switching.
+
+# v4.2.0-beta
+- [#1202](https://github.com/xmrig/xmrig/issues/1202) Fixed algorithm verification in donate strategy.
+- Added per pool option `coin` with single possible value `monero` for pools without algorithm negotiation, for upcoming Monero fork.
+- Added config option `cpu/max-threads-hint` and command line option `--cpu-max-threads-hint`.
+
+# v4.1.0-beta
+- **OpenCL backend disabled by default.**.
+- [#1183](https://github.com/xmrig/xmrig/issues/1183) Fixed compatibility with systemd.
+- [#1185](https://github.com/xmrig/xmrig/pull/1185) Added JIT compiler for RandomX on ARMv8.
+- Improved API endpoint `GET /2/backends` and added support for this endpoint to [workers.xmrig.info](http://workers.xmrig.info).
+- Added command line option `--no-cpu` to disable CPU backend.
+- Added OpenCL specific command line options: `--opencl`, `--opencl-devices`, `--opencl-platform`, `--opencl-loader` and `--opencl-no-cache`.
+- Removed command line option `--http-enabled`, HTTP API enabled automatically if any other `--http-*` option provided.
+
+# v4.0.1-beta
+- [#1177](https://github.com/xmrig/xmrig/issues/1177) Fixed compatibility with old AMD drivers.
+- [#1180](https://github.com/xmrig/xmrig/issues/1180) Fixed possible duplicated shares after algorithm switching.
+- Added support for case if not all backend threads successfully started.
+- Fixed wrong config file permissions after write (only gcc builds on recent Windows 10 affected).
+
 # v4.0.0-beta
 - [#1172](https://github.com/xmrig/xmrig/issues/1172) **Added OpenCL mining backend.**
   - [#268](https://github.com/xmrig/xmrig-amd/pull/268) [#270](https://github.com/xmrig/xmrig-amd/pull/270) [#271](https://github.com/xmrig/xmrig-amd/pull/271) [#273](https://github.com/xmrig/xmrig-amd/pull/273) [#274](https://github.com/xmrig/xmrig-amd/pull/274) [#1171](https://github.com/xmrig/xmrig/pull/1171) Added RandomX support for OpenCL, thanks [@SChernykh](https://github.com/SChernykh).
 - Algorithm `cn/wow` removed, as no longer alive. 
 
+# v3.2.0
+- Added per pool option `coin` with single possible value `monero` for pools without algorithm negotiation, for upcoming Monero fork.
+- [#1183](https://github.com/xmrig/xmrig/issues/1183) Fixed compatibility with systemd.
+
+# v3.1.3
+- [#1180](https://github.com/xmrig/xmrig/issues/1180) Fixed possible duplicated shares after algorithm switching.
+- Fixed wrong config file permissions after write (only gcc builds on recent Windows 10 affected).
+
 # v3.1.2
 - Many RandomX optimizations and fixes.
   - [#1132](https://github.com/xmrig/xmrig/issues/1132) Fixed build on CentOS 7.

CMakeLists.txt

@@ -71,6 +71,7 @@ set(HEADERS_CRYPTO
     src/crypto/cn/skein_port.h
     src/crypto/cn/soft_aes.h
     src/crypto/common/Algorithm.h
+    src/crypto/common/Coin.h
     src/crypto/common/keccak.h
     src/crypto/common/Nonce.h
     src/crypto/common/portable/mm_malloc.h
@@ -108,6 +109,7 @@ set(SOURCES_CRYPTO
     src/crypto/cn/CnCtx.cpp
     src/crypto/cn/CnHash.cpp
     src/crypto/common/Algorithm.cpp
+    src/crypto/common/Coin.cpp
     src/crypto/common/keccak.cpp
     src/crypto/common/Nonce.cpp
     src/crypto/common/VirtualMemory.cpp
@@ -143,7 +145,7 @@ else()
     endif()
 endif()
 
-if (CMAKE_SYSTEM_NAME MATCHES "Linux")
+if (CMAKE_SYSTEM_NAME MATCHES "Linux" OR CMAKE_SYSTEM_NAME MATCHES "Android")
     EXECUTE_PROCESS(COMMAND uname -o COMMAND tr -d '\n' OUTPUT_VARIABLE OPERATING_SYSTEM)
     if (OPERATING_SYSTEM MATCHES "Android")
         set(EXTRA_LIBS ${EXTRA_LIBS} log)

README.md

@@ -1,5 +1,7 @@
 # XMRig
 
+**:warning: [Monero will change PoW algorithm to RandomX on November 30.](https://github.com/xmrig/xmrig/issues/1204)**
+
 [![Github All Releases](https://img.shields.io/github/downloads/xmrig/xmrig/total.svg)](https://github.com/xmrig/xmrig/releases)
 [![GitHub release](https://img.shields.io/github/release/xmrig/xmrig/all.svg)](https://github.com/xmrig/xmrig/releases)
 [![GitHub Release Date](https://img.shields.io/github/release-date-pre/xmrig/xmrig.svg)](https://github.com/xmrig/xmrig/releases)
@@ -26,7 +28,7 @@ XMRig is a high performance RandomX and CryptoNight CPU miner, with official sup
   * Clone with `git clone https://github.com/xmrig/xmrig.git` :hammer: [Build instructions](https://github.com/xmrig/xmrig/wiki/Build).
 
 ## Usage
-Preferend way to configure miner is [JSON config file](src/config.json) as more flexible and human frendly, command line interface not cover all features, for example mining profiles for different algorithms. Most impotant options can be changed in runtime without miner restart by editing config or via API.
+The preferred way to configure the miner is the [JSON config file](src/config.json) as it is more flexible and human friendly. The command line interface does not cover all features, such as mining profiles for different algorithms. Important options can be changed during runtime without miner restart by editing the config file or executing API calls.
 
 ### Options
 ```

cmake/randomx.cmake

@@ -51,6 +51,13 @@ if (WITH_RANDOMX)
             )
         # cheat because cmake and ccache hate each other
         set_property(SOURCE src/crypto/randomx/jit_compiler_x86_static.S PROPERTY LANGUAGE C)
+    elseif (XMRIG_ARM AND CMAKE_SIZEOF_VOID_P EQUAL 8)
+        list(APPEND SOURCES_CRYPTO
+             src/crypto/randomx/jit_compiler_a64_static.S
+             src/crypto/randomx/jit_compiler_a64.cpp
+            )
+        # cheat because cmake and ccache hate each other
+        set_property(SOURCE src/crypto/randomx/jit_compiler_a64_static.S PROPERTY LANGUAGE C)
     endif()
 
     if (CMAKE_CXX_COMPILER_ID MATCHES Clang)

doc/ALGORITHMS.md

@@ -1,34 +1,20 @@
 # Algorithms
 
-Since version 3 mining [algorithm](#algorithm-names) should specified for each pool separately (`algo` option), earlier versions was use one global `algo` option and per pool `variant` option (this option was removed in v3). If your pool support [mining algorithm negotiation](https://github.com/xmrig/xmrig-proxy/issues/168) you may not specify this option at all.
- 
-#### Example
-```json
-{
-  "pools": [
-    {
-      "url": "...",
-      "algo": "cn/r",
-      ...
-    }
- ],
- ...
-}
-```
+Algorithm can be defined in 3 ways:
 
-#### Pools with mining algorithm negotiation support.
+1. By pool, using algorithm negotiation, in this case no need specify algorithm on miner side.
+2. Per pool `coin` option, currently only usable value for this option is `monero`.
+3. Per pool `algo` option.
 
- * [www.hashvault.pro](https://www.hashvault.pro/)
- * [moneroocean.stream](https://moneroocean.stream)
- 
- ## Algorithm names
+Option `coin` useful for pools without algorithm negotiation support or daemon to allow automatically switch algorithm in next hard fork.
+
+## Algorithm names
 
 | Name | Memory | Version | Notes |
 |------|--------|---------|-------|
+| `rx/0` | 2 MB | 3.2.0+ | RandomX (Monero). |
 | `argon2/chukwa` | 512 KB | 3.1.0+ | Argon2id (Chukwa). |
 | `argon2/wrkz` | 256 KB | 3.1.0+ | Argon2id (WRKZ) |
-| `rx/test` | 2 MB | 3.0.0+ | RandomX (reference configuration). |
-| `rx/0` | 2 MB | 3.0.0+ | RandomX (reference configuration), reserved for future use. |
 | `rx/wow` | 1 MB | 3.0.0+ | RandomWOW. |
 | `rx/loki` | 2 MB | 3.0.0+ | RandomXL. |
 | `cn/fast` | 2 MB | 3.0.0+ | CryptoNight variant 1 with half iterations. |
@@ -50,3 +36,21 @@ Since version 3 mining [algorithm](#algorithm-names) should specified for each p
 | `cn-lite/1` | 1 MB | 2.5.0+ | CryptoNight-Lite variant 1. |
 | `cn-lite/0` | 1 MB | 0.8.0+ | CryptoNight-Lite variant 0. |
 | `cn/0` | 2 MB | 0.5.0+ | CryptoNight (original). |
+
+## Migration to v3
+Since version 3 mining [algorithm](#algorithm-names) should specified for each pool separately (`algo` option), earlier versions was use one global `algo` option and per pool `variant` option (this option was removed in v3). If your pool support [mining algorithm negotiation](https://github.com/xmrig/xmrig-proxy/issues/168) you may not specify this option at all.
+ 
+#### Example
+```json
+{
+  "pools": [
+    {
+      "url": "...",
+      "algo": "cn/r",
+      "coin": null
+      ...
+    }
+ ],
+ ...
+}
+```

doc/CPU.md

@@ -94,3 +94,6 @@ Enable/configure or disable ASM optimizations. Possible values: `true`, `false`,
 
 #### `argon2-impl` (since v3.1.0)
 Allow override automatically detected Argon2 implementation, this option added mostly for debug purposes, default value `null` means autodetect. Other possible values: `"x86_64"`, `"SSE2"`, `"SSSE3"`, `"XOP"`, `"AVX2"`, `"AVX-512F"`. Manual selection has no safe guards, if you CPU not support required instuctions, miner will crash.
+
+#### `max-threads-hint` (since v4.2.0)
+Maximum CPU threads count (in percentage) hint for autoconfig. [CPU_MAX_USAGE.md](CPU_MAX_USAGE.md)

doc/CPU_MAX_USAGE.md

@@ -0,0 +1,26 @@
+# Maximum CPU usage
+
+Please read this document carefully, `max-threads-hint` (was known as `max-cpu-usage`) option is most confusing option in the miner with many myth and legends.
+This option is just hint for automatic configuration and can't precise define CPU usage.
+
+### Option definition
+#### Config file:
+```json
+{
+    ...
+    "cpu": {
+        "max-threads-hint": 100,
+        ...
+    },
+    ...
+}
+```
+
+#### Command line
+`--cpu-max-threads-hint 100`
+
+### Known issues and usage
+
+* This option has no effect if miner already generated CPU configuration, to prevent config generation use `"autosave":false,`.
+* Only threads count can be changed, for 1 core CPU this option has no effect, for 2 core CPU only 2 values possible 50% and 100%, for 4 cores: 25%, 50%, 75%, 100%. etc. 
+* You CPU may limited by other factors, eg cache.

doc/topology/AMD_Opteron_6344_x2_N4_win7_2_0_4_bug.xml

@@ -0,0 +1,262 @@
+<?xml version="1.0" encoding="UTF-8"?>
+<!DOCTYPE topology SYSTEM "hwloc2.dtd">
+<topology version="2.0">
+  <object type="Machine" os_index="0" cpuset="0x00ffffff" complete_cpuset="0x00ffffff" allowed_cpuset="0x00ffffff" nodeset="0x0000000f" complete_nodeset="0x0000000f" allowed_nodeset="0x0000000f" gp_index="1">
+    <info name="Backend" value="Windows"/>
+    <info name="hwlocVersion" value="2.0.4"/>
+    <object type="Package" cpuset="0x00000fff" complete_cpuset="0x00000fff" nodeset="0x00000003" complete_nodeset="0x00000003" gp_index="36">
+      <info name="CPUVendor" value="AuthenticAMD"/>
+      <info name="CPUFamilyNumber" value="21"/>
+      <info name="CPUModelNumber" value="2"/>
+      <info name="CPUModel" value="AMD Opteron(tm) Processor 6344                 "/>
+      <info name="CPUStepping" value="0"/>
+      <object type="L3Cache" cpuset="0x0000003f" complete_cpuset="0x0000003f" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="20" cache_size="12582912" depth="3" cache_linesize="64" cache_associativity="1" cache_type="0">
+        <info name="Inclusive" value="0"/>
+        <object type="NUMANode" os_index="0" cpuset="0x0000003f" complete_cpuset="0x0000003f" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="80" local_memory="7009357824">
+          <page_type size="4096" count="0"/>
+        </object>
+        <object type="L2Cache" cpuset="0x00000001" complete_cpuset="0x00000001" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="4" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000001" complete_cpuset="0x00000001" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="3" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00000001" complete_cpuset="0x00000001" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="2">
+              <object type="PU" os_index="0" cpuset="0x00000001" complete_cpuset="0x00000001" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="85"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000002" complete_cpuset="0x00000002" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="7" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000002" complete_cpuset="0x00000002" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="6" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00000002" complete_cpuset="0x00000002" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="5">
+              <object type="PU" os_index="1" cpuset="0x00000002" complete_cpuset="0x00000002" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="86"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000004" complete_cpuset="0x00000004" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="10" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000004" complete_cpuset="0x00000004" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="9" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00000004" complete_cpuset="0x00000004" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="8">
+              <object type="PU" os_index="2" cpuset="0x00000004" complete_cpuset="0x00000004" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="87"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000008" complete_cpuset="0x00000008" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="13" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000008" complete_cpuset="0x00000008" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="12" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00000008" complete_cpuset="0x00000008" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="11">
+              <object type="PU" os_index="3" cpuset="0x00000008" complete_cpuset="0x00000008" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="88"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000010" complete_cpuset="0x00000010" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="16" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000010" complete_cpuset="0x00000010" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="15" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00000010" complete_cpuset="0x00000010" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="14">
+              <object type="PU" os_index="4" cpuset="0x00000010" complete_cpuset="0x00000010" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="89"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000020" complete_cpuset="0x00000020" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="19" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000020" complete_cpuset="0x00000020" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="18" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00000020" complete_cpuset="0x00000020" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="17">
+              <object type="PU" os_index="5" cpuset="0x00000020" complete_cpuset="0x00000020" nodeset="0x00000001" complete_nodeset="0x00000001" gp_index="90"/>
+            </object>
+          </object>
+        </object>
+      </object>
+      <object type="L3Cache" cpuset="0x00000fc0" complete_cpuset="0x00000fc0" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="40" cache_size="12582912" depth="3" cache_linesize="64" cache_associativity="1" cache_type="0">
+        <info name="Inclusive" value="0"/>
+        <object type="NUMANode" os_index="1" cpuset="0x00000fc0" complete_cpuset="0x00000fc0" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="81" local_memory="8018194432">
+          <page_type size="4096" count="0"/>
+        </object>
+        <object type="L2Cache" cpuset="0x00000040" complete_cpuset="0x00000040" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="23" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000040" complete_cpuset="0x00000040" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="22" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00000040" complete_cpuset="0x00000040" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="21">
+              <object type="PU" os_index="6" cpuset="0x00000040" complete_cpuset="0x00000040" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="91"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000080" complete_cpuset="0x00000080" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="26" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000080" complete_cpuset="0x00000080" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="25" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00000080" complete_cpuset="0x00000080" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="24">
+              <object type="PU" os_index="7" cpuset="0x00000080" complete_cpuset="0x00000080" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="92"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000100" complete_cpuset="0x00000100" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="29" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000100" complete_cpuset="0x00000100" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="28" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00000100" complete_cpuset="0x00000100" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="27">
+              <object type="PU" os_index="8" cpuset="0x00000100" complete_cpuset="0x00000100" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="93"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000200" complete_cpuset="0x00000200" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="32" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000200" complete_cpuset="0x00000200" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="31" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00000200" complete_cpuset="0x00000200" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="30">
+              <object type="PU" os_index="9" cpuset="0x00000200" complete_cpuset="0x00000200" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="94"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000400" complete_cpuset="0x00000400" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="35" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000400" complete_cpuset="0x00000400" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="34" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00000400" complete_cpuset="0x00000400" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="33">
+              <object type="PU" os_index="10" cpuset="0x00000400" complete_cpuset="0x00000400" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="95"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000800" complete_cpuset="0x00000800" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="39" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000800" complete_cpuset="0x00000800" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="38" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00000800" complete_cpuset="0x00000800" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="37">
+              <object type="PU" os_index="11" cpuset="0x00000800" complete_cpuset="0x00000800" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="96"/>
+            </object>
+          </object>
+        </object>
+      </object>
+    </object>
+    <object type="Package" cpuset="0x00fff000" complete_cpuset="0x00fff000" nodeset="0x0000000c" complete_nodeset="0x0000000c" gp_index="75">
+      <info name="CPUVendor" value="AuthenticAMD"/>
+      <info name="CPUFamilyNumber" value="21"/>
+      <info name="CPUModelNumber" value="2"/>
+      <info name="CPUModel" value="AMD Opteron(tm) Processor 6344                 "/>
+      <info name="CPUStepping" value="0"/>
+      <object type="L3Cache" cpuset="0x0003f000" complete_cpuset="0x0003f000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="59" cache_size="12582912" depth="3" cache_linesize="64" cache_associativity="1" cache_type="0">
+        <info name="Inclusive" value="0"/>
+        <object type="NUMANode" os_index="2" cpuset="0x0003f000" complete_cpuset="0x0003f000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="82" local_memory="8035020800">
+          <page_type size="4096" count="0"/>
+        </object>
+        <object type="L2Cache" cpuset="0x00001000" complete_cpuset="0x00001000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="43" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00001000" complete_cpuset="0x00001000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="42" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00001000" complete_cpuset="0x00001000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="41">
+              <object type="PU" os_index="12" cpuset="0x00001000" complete_cpuset="0x00001000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="97"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00002000" complete_cpuset="0x00002000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="46" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00002000" complete_cpuset="0x00002000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="45" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00002000" complete_cpuset="0x00002000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="44">
+              <object type="PU" os_index="13" cpuset="0x00002000" complete_cpuset="0x00002000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="98"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00004000" complete_cpuset="0x00004000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="49" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00004000" complete_cpuset="0x00004000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="48" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00004000" complete_cpuset="0x00004000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="47">
+              <object type="PU" os_index="14" cpuset="0x00004000" complete_cpuset="0x00004000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="99"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00008000" complete_cpuset="0x00008000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="52" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00008000" complete_cpuset="0x00008000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="51" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00008000" complete_cpuset="0x00008000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="50">
+              <object type="PU" os_index="15" cpuset="0x00008000" complete_cpuset="0x00008000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="100"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00010000" complete_cpuset="0x00010000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="55" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00010000" complete_cpuset="0x00010000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="54" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00010000" complete_cpuset="0x00010000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="53">
+              <object type="PU" os_index="16" cpuset="0x00010000" complete_cpuset="0x00010000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="101"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00020000" complete_cpuset="0x00020000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="58" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00020000" complete_cpuset="0x00020000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="57" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00020000" complete_cpuset="0x00020000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="56">
+              <object type="PU" os_index="17" cpuset="0x00020000" complete_cpuset="0x00020000" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="102"/>
+            </object>
+          </object>
+        </object>
+      </object>
+      <object type="L3Cache" cpuset="0x00fc0000" complete_cpuset="0x00fc0000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="79" cache_size="12582912" depth="3" cache_linesize="64" cache_associativity="1" cache_type="0">
+        <info name="Inclusive" value="0"/>
+        <object type="NUMANode" os_index="3" cpuset="0x00fc0000" complete_cpuset="0x00fc0000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="83" local_memory="8097337344">
+          <page_type size="4096" count="0"/>
+        </object>
+        <object type="L2Cache" cpuset="0x00040000" complete_cpuset="0x00040000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="62" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00040000" complete_cpuset="0x00040000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="61" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00040000" complete_cpuset="0x00040000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="60">
+              <object type="PU" os_index="18" cpuset="0x00040000" complete_cpuset="0x00040000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="103"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00080000" complete_cpuset="0x00080000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="65" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00080000" complete_cpuset="0x00080000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="64" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00080000" complete_cpuset="0x00080000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="63">
+              <object type="PU" os_index="19" cpuset="0x00080000" complete_cpuset="0x00080000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="104"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00100000" complete_cpuset="0x00100000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="68" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00100000" complete_cpuset="0x00100000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="67" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00100000" complete_cpuset="0x00100000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="66">
+              <object type="PU" os_index="20" cpuset="0x00100000" complete_cpuset="0x00100000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="105"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00200000" complete_cpuset="0x00200000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="71" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00200000" complete_cpuset="0x00200000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="70" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00200000" complete_cpuset="0x00200000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="69">
+              <object type="PU" os_index="21" cpuset="0x00200000" complete_cpuset="0x00200000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="106"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00400000" complete_cpuset="0x00400000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="74" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00400000" complete_cpuset="0x00400000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="73" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00400000" complete_cpuset="0x00400000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="72">
+              <object type="PU" os_index="22" cpuset="0x00400000" complete_cpuset="0x00400000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="107"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00800000" complete_cpuset="0x00800000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="78" cache_size="2097152" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00800000" complete_cpuset="0x00800000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="77" cache_size="16384" depth="1" cache_linesize="64" cache_associativity="4" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" cpuset="0x00800000" complete_cpuset="0x00800000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="76">
+              <object type="PU" os_index="23" cpuset="0x00800000" complete_cpuset="0x00800000" nodeset="0x00000008" complete_nodeset="0x00000008" gp_index="108"/>
+            </object>
+          </object>
+        </object>
+      </object>
+    </object>
+  </object>
+</topology>

doc/topology/AMD_Opteron_6348_x4_N8_linux_2_0_4.xml

@@ -0,0 +1,399 @@
+<?xml version="1.0" encoding="UTF-8"?>
+<!DOCTYPE topology SYSTEM "hwloc2.dtd">
+<topology version="2.0">
+  <object type="Machine" os_index="0" cpuset="0xffffffff" complete_cpuset="0xffffffff" allowed_cpuset="0xffffffff" nodeset="0x00000066" complete_nodeset="0x000000ff" allowed_nodeset="0x00000066" gp_index="1">
+    <info name="DMIProductName" value="H8QG6"/>
+    <info name="DMIProductVersion" value="1234567890"/>
+    <info name="DMIProductSerial" value="1234567890"/>
+    <info name="DMIProductUUID" value="0"/>
+    <info name="DMIBoardVendor" value="Supermicro"/>
+    <info name="DMIBoardName" value="H8QG6"/>
+    <info name="DMIBoardVersion" value="1234567890"/>
+    <info name="DMIBoardSerial" value="0"/>
+    <info name="DMIBoardAssetTag" value="1234567890"/>
+    <info name="DMIChassisVendor" value="Supermicro"/>
+    <info name="DMIChassisType" value="3"/>
+    <info name="DMIChassisVersion" value="1234567890"/>
+    <info name="DMIChassisSerial" value="1234567890."/>
+    <info name="DMIChassisAssetTag" value="1234567890"/>
+    <info name="DMIBIOSVendor" value="American Megatrends Inc."/>
+    <info name="DMIBIOSVersion" value="080016 "/>
+    <info name="DMIBIOSDate" value="10/11/2010"/>
+    <info name="DMISysVendor" value="Supermicro"/>
+    <info name="Backend" value="Linux"/>
+    <info name="LinuxCgroup" value="/"/>
+    <info name="OSName" value="Linux"/>
+    <info name="OSRelease" value="4.15.0-20-generic"/>
+    <info name="OSVersion" value="#21-Ubuntu SMP Tue Apr 24 06:16:15 UTC 2018"/>
+    <info name="HostName" value="host"/>
+    <info name="Architecture" value="x86_64"/>
+    <info name="hwlocVersion" value="2.0.4"/>
+    <info name="ProcessName" value="xmrig"/>
+    <object type="Package" os_index="0" cpuset="0x000000ff" complete_cpuset="0x000000ff" nodeset="0x00000002" complete_nodeset="0x00000003" gp_index="2">
+      <info name="CPUVendor" value="AuthenticAMD"/>
+      <info name="CPUFamilyNumber" value="16"/>
+      <info name="CPUModelNumber" value="9"/>
+      <info name="CPUModel" value="AMD Opteron(tm) Processor 6128"/>
+      <info name="CPUStepping" value="1"/>
+      <object type="L3Cache" cpuset="0x0000000f" complete_cpuset="0x0000000f" nodeset="0x0" complete_nodeset="0x00000001" gp_index="7" cache_size="5240832" depth="3" cache_linesize="64" cache_associativity="48" cache_type="0">
+        <info name="Inclusive" value="0"/>
+        <object type="L2Cache" cpuset="0x00000001" complete_cpuset="0x00000001" nodeset="0x0" complete_nodeset="0x00000001" gp_index="6" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000001" complete_cpuset="0x00000001" nodeset="0x0" complete_nodeset="0x00000001" gp_index="5" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="0" cpuset="0x00000001" complete_cpuset="0x00000001" nodeset="0x0" complete_nodeset="0x00000001" gp_index="3">
+              <object type="PU" os_index="0" cpuset="0x00000001" complete_cpuset="0x00000001" nodeset="0x0" complete_nodeset="0x00000001" gp_index="4"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000002" complete_cpuset="0x00000002" nodeset="0x0" complete_nodeset="0x00000001" gp_index="11" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000002" complete_cpuset="0x00000002" nodeset="0x0" complete_nodeset="0x00000001" gp_index="10" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="1" cpuset="0x00000002" complete_cpuset="0x00000002" nodeset="0x0" complete_nodeset="0x00000001" gp_index="8">
+              <object type="PU" os_index="1" cpuset="0x00000002" complete_cpuset="0x00000002" nodeset="0x0" complete_nodeset="0x00000001" gp_index="9"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000004" complete_cpuset="0x00000004" nodeset="0x0" complete_nodeset="0x00000001" gp_index="15" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000004" complete_cpuset="0x00000004" nodeset="0x0" complete_nodeset="0x00000001" gp_index="14" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="2" cpuset="0x00000004" complete_cpuset="0x00000004" nodeset="0x0" complete_nodeset="0x00000001" gp_index="12">
+              <object type="PU" os_index="2" cpuset="0x00000004" complete_cpuset="0x00000004" nodeset="0x0" complete_nodeset="0x00000001" gp_index="13"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000008" complete_cpuset="0x00000008" nodeset="0x0" complete_nodeset="0x00000001" gp_index="19" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000008" complete_cpuset="0x00000008" nodeset="0x0" complete_nodeset="0x00000001" gp_index="18" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="3" cpuset="0x00000008" complete_cpuset="0x00000008" nodeset="0x0" complete_nodeset="0x00000001" gp_index="16">
+              <object type="PU" os_index="3" cpuset="0x00000008" complete_cpuset="0x00000008" nodeset="0x0" complete_nodeset="0x00000001" gp_index="17"/>
+            </object>
+          </object>
+        </object>
+      </object>
+      <object type="L3Cache" cpuset="0x000000f0" complete_cpuset="0x000000f0" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="24" cache_size="5240832" depth="3" cache_linesize="64" cache_associativity="48" cache_type="0">
+        <info name="Inclusive" value="0"/>
+        <object type="NUMANode" os_index="1" cpuset="0x000000f0" complete_cpuset="0x000000f0" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="143" local_memory="4156817408">
+          <page_type size="4096" count="854592"/>
+          <page_type size="2097152" count="313"/>
+          <page_type size="1073741824" count="0"/>
+        </object>
+        <object type="L2Cache" cpuset="0x00000010" complete_cpuset="0x00000010" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="23" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000010" complete_cpuset="0x00000010" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="22" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="0" cpuset="0x00000010" complete_cpuset="0x00000010" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="20">
+              <object type="PU" os_index="4" cpuset="0x00000010" complete_cpuset="0x00000010" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="21"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000020" complete_cpuset="0x00000020" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="28" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000020" complete_cpuset="0x00000020" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="27" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="1" cpuset="0x00000020" complete_cpuset="0x00000020" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="25">
+              <object type="PU" os_index="5" cpuset="0x00000020" complete_cpuset="0x00000020" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="26"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000040" complete_cpuset="0x00000040" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="32" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000040" complete_cpuset="0x00000040" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="31" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="2" cpuset="0x00000040" complete_cpuset="0x00000040" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="29">
+              <object type="PU" os_index="6" cpuset="0x00000040" complete_cpuset="0x00000040" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="30"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000080" complete_cpuset="0x00000080" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="36" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000080" complete_cpuset="0x00000080" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="35" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="3" cpuset="0x00000080" complete_cpuset="0x00000080" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="33">
+              <object type="PU" os_index="7" cpuset="0x00000080" complete_cpuset="0x00000080" nodeset="0x00000002" complete_nodeset="0x00000002" gp_index="34"/>
+            </object>
+          </object>
+        </object>
+      </object>
+    </object>
+    <object type="Package" os_index="1" cpuset="0x0000ff00" complete_cpuset="0x0000ff00" nodeset="0x00000004" complete_nodeset="0x0000000c" gp_index="37">
+      <info name="CPUVendor" value="AuthenticAMD"/>
+      <info name="CPUFamilyNumber" value="16"/>
+      <info name="CPUModelNumber" value="9"/>
+      <info name="CPUModel" value="AMD Opteron(tm) Processor 6128"/>
+      <info name="CPUStepping" value="1"/>
+      <object type="L3Cache" cpuset="0x00000f00" complete_cpuset="0x00000f00" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="42" cache_size="5240832" depth="3" cache_linesize="64" cache_associativity="48" cache_type="0">
+        <info name="Inclusive" value="0"/>
+        <object type="NUMANode" os_index="2" cpuset="0x00000f00" complete_cpuset="0x00000f00" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="144" local_memory="4204060672">
+          <page_type size="4096" count="866126"/>
+          <page_type size="2097152" count="313"/>
+          <page_type size="1073741824" count="0"/>
+        </object>
+        <object type="L2Cache" cpuset="0x00000100" complete_cpuset="0x00000100" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="41" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000100" complete_cpuset="0x00000100" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="40" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="0" cpuset="0x00000100" complete_cpuset="0x00000100" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="38">
+              <object type="PU" os_index="8" cpuset="0x00000100" complete_cpuset="0x00000100" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="39"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000200" complete_cpuset="0x00000200" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="46" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000200" complete_cpuset="0x00000200" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="45" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="1" cpuset="0x00000200" complete_cpuset="0x00000200" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="43">
+              <object type="PU" os_index="9" cpuset="0x00000200" complete_cpuset="0x00000200" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="44"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000400" complete_cpuset="0x00000400" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="50" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000400" complete_cpuset="0x00000400" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="49" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="2" cpuset="0x00000400" complete_cpuset="0x00000400" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="47">
+              <object type="PU" os_index="10" cpuset="0x00000400" complete_cpuset="0x00000400" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="48"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00000800" complete_cpuset="0x00000800" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="54" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00000800" complete_cpuset="0x00000800" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="53" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="3" cpuset="0x00000800" complete_cpuset="0x00000800" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="51">
+              <object type="PU" os_index="11" cpuset="0x00000800" complete_cpuset="0x00000800" nodeset="0x00000004" complete_nodeset="0x00000004" gp_index="52"/>
+            </object>
+          </object>
+        </object>
+      </object>
+      <object type="L3Cache" cpuset="0x0000f000" complete_cpuset="0x0000f000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="59" cache_size="5240832" depth="3" cache_linesize="64" cache_associativity="48" cache_type="0">
+        <info name="Inclusive" value="0"/>
+        <object type="L2Cache" cpuset="0x00001000" complete_cpuset="0x00001000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="58" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00001000" complete_cpuset="0x00001000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="57" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="0" cpuset="0x00001000" complete_cpuset="0x00001000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="55">
+              <object type="PU" os_index="12" cpuset="0x00001000" complete_cpuset="0x00001000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="56"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00002000" complete_cpuset="0x00002000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="63" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00002000" complete_cpuset="0x00002000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="62" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="1" cpuset="0x00002000" complete_cpuset="0x00002000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="60">
+              <object type="PU" os_index="13" cpuset="0x00002000" complete_cpuset="0x00002000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="61"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00004000" complete_cpuset="0x00004000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="67" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00004000" complete_cpuset="0x00004000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="66" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="2" cpuset="0x00004000" complete_cpuset="0x00004000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="64">
+              <object type="PU" os_index="14" cpuset="0x00004000" complete_cpuset="0x00004000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="65"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00008000" complete_cpuset="0x00008000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="71" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00008000" complete_cpuset="0x00008000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="70" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="3" cpuset="0x00008000" complete_cpuset="0x00008000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="68">
+              <object type="PU" os_index="15" cpuset="0x00008000" complete_cpuset="0x00008000" nodeset="0x0" complete_nodeset="0x00000008" gp_index="69"/>
+            </object>
+          </object>
+        </object>
+      </object>
+    </object>
+    <object type="Package" os_index="2" cpuset="0x00ff0000" complete_cpuset="0x00ff0000" nodeset="0x00000020" complete_nodeset="0x00000030" gp_index="72">
+      <info name="CPUVendor" value="AuthenticAMD"/>
+      <info name="CPUFamilyNumber" value="16"/>
+      <info name="CPUModelNumber" value="9"/>
+      <info name="CPUModel" value="AMD Opteron(tm) Processor 6128"/>
+      <info name="CPUStepping" value="1"/>
+      <object type="L3Cache" cpuset="0x000f0000" complete_cpuset="0x000f0000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="77" cache_size="5240832" depth="3" cache_linesize="64" cache_associativity="48" cache_type="0">
+        <info name="Inclusive" value="0"/>
+        <object type="L2Cache" cpuset="0x00010000" complete_cpuset="0x00010000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="76" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00010000" complete_cpuset="0x00010000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="75" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="0" cpuset="0x00010000" complete_cpuset="0x00010000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="73">
+              <object type="PU" os_index="16" cpuset="0x00010000" complete_cpuset="0x00010000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="74"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00020000" complete_cpuset="0x00020000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="81" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00020000" complete_cpuset="0x00020000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="80" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="1" cpuset="0x00020000" complete_cpuset="0x00020000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="78">
+              <object type="PU" os_index="17" cpuset="0x00020000" complete_cpuset="0x00020000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="79"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00040000" complete_cpuset="0x00040000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="85" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00040000" complete_cpuset="0x00040000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="84" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="2" cpuset="0x00040000" complete_cpuset="0x00040000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="82">
+              <object type="PU" os_index="18" cpuset="0x00040000" complete_cpuset="0x00040000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="83"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00080000" complete_cpuset="0x00080000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="89" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00080000" complete_cpuset="0x00080000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="88" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="3" cpuset="0x00080000" complete_cpuset="0x00080000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="86">
+              <object type="PU" os_index="19" cpuset="0x00080000" complete_cpuset="0x00080000" nodeset="0x0" complete_nodeset="0x00000010" gp_index="87"/>
+            </object>
+          </object>
+        </object>
+      </object>
+      <object type="L3Cache" cpuset="0x00f00000" complete_cpuset="0x00f00000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="94" cache_size="5240832" depth="3" cache_linesize="64" cache_associativity="48" cache_type="0">
+        <info name="Inclusive" value="0"/>
+        <object type="NUMANode" os_index="5" cpuset="0x00f00000" complete_cpuset="0x00f00000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="147" local_memory="4226170880">
+          <page_type size="4096" count="872036"/>
+          <page_type size="2097152" count="312"/>
+          <page_type size="1073741824" count="0"/>
+        </object>
+        <object type="L2Cache" cpuset="0x00100000" complete_cpuset="0x00100000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="93" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00100000" complete_cpuset="0x00100000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="92" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="0" cpuset="0x00100000" complete_cpuset="0x00100000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="90">
+              <object type="PU" os_index="20" cpuset="0x00100000" complete_cpuset="0x00100000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="91"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00200000" complete_cpuset="0x00200000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="98" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00200000" complete_cpuset="0x00200000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="97" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="1" cpuset="0x00200000" complete_cpuset="0x00200000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="95">
+              <object type="PU" os_index="21" cpuset="0x00200000" complete_cpuset="0x00200000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="96"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00400000" complete_cpuset="0x00400000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="102" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00400000" complete_cpuset="0x00400000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="101" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="2" cpuset="0x00400000" complete_cpuset="0x00400000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="99">
+              <object type="PU" os_index="22" cpuset="0x00400000" complete_cpuset="0x00400000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="100"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x00800000" complete_cpuset="0x00800000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="106" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x00800000" complete_cpuset="0x00800000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="105" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="3" cpuset="0x00800000" complete_cpuset="0x00800000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="103">
+              <object type="PU" os_index="23" cpuset="0x00800000" complete_cpuset="0x00800000" nodeset="0x00000020" complete_nodeset="0x00000020" gp_index="104"/>
+            </object>
+          </object>
+        </object>
+      </object>
+    </object>
+    <object type="Package" os_index="3" cpuset="0xff000000" complete_cpuset="0xff000000" nodeset="0x00000040" complete_nodeset="0x000000c0" gp_index="107">
+      <info name="CPUVendor" value="AuthenticAMD"/>
+      <info name="CPUFamilyNumber" value="16"/>
+      <info name="CPUModelNumber" value="9"/>
+      <info name="CPUModel" value="AMD Opteron(tm) Processor 6128"/>
+      <info name="CPUStepping" value="1"/>
+      <object type="L3Cache" cpuset="0x0f000000" complete_cpuset="0x0f000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="112" cache_size="5240832" depth="3" cache_linesize="64" cache_associativity="48" cache_type="0">
+        <info name="Inclusive" value="0"/>
+        <object type="NUMANode" os_index="6" cpuset="0x0f000000" complete_cpuset="0x0f000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="148" local_memory="4221870080">
+          <page_type size="4096" count="870986"/>
+          <page_type size="2097152" count="312"/>
+          <page_type size="1073741824" count="0"/>
+        </object>
+        <object type="L2Cache" cpuset="0x01000000" complete_cpuset="0x01000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="111" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x01000000" complete_cpuset="0x01000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="110" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="0" cpuset="0x01000000" complete_cpuset="0x01000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="108">
+              <object type="PU" os_index="24" cpuset="0x01000000" complete_cpuset="0x01000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="109"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x02000000" complete_cpuset="0x02000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="116" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x02000000" complete_cpuset="0x02000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="115" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="1" cpuset="0x02000000" complete_cpuset="0x02000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="113">
+              <object type="PU" os_index="25" cpuset="0x02000000" complete_cpuset="0x02000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="114"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x04000000" complete_cpuset="0x04000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="120" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x04000000" complete_cpuset="0x04000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="119" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="2" cpuset="0x04000000" complete_cpuset="0x04000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="117">
+              <object type="PU" os_index="26" cpuset="0x04000000" complete_cpuset="0x04000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="118"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x08000000" complete_cpuset="0x08000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="124" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x08000000" complete_cpuset="0x08000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="123" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="3" cpuset="0x08000000" complete_cpuset="0x08000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="121">
+              <object type="PU" os_index="27" cpuset="0x08000000" complete_cpuset="0x08000000" nodeset="0x00000040" complete_nodeset="0x00000040" gp_index="122"/>
+            </object>
+          </object>
+        </object>
+      </object>
+      <object type="L3Cache" cpuset="0xf0000000" complete_cpuset="0xf0000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="129" cache_size="5240832" depth="3" cache_linesize="64" cache_associativity="48" cache_type="0">
+        <info name="Inclusive" value="0"/>
+        <object type="L2Cache" cpuset="0x10000000" complete_cpuset="0x10000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="128" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x10000000" complete_cpuset="0x10000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="127" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="0" cpuset="0x10000000" complete_cpuset="0x10000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="125">
+              <object type="PU" os_index="28" cpuset="0x10000000" complete_cpuset="0x10000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="126"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x20000000" complete_cpuset="0x20000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="133" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x20000000" complete_cpuset="0x20000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="132" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="1" cpuset="0x20000000" complete_cpuset="0x20000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="130">
+              <object type="PU" os_index="29" cpuset="0x20000000" complete_cpuset="0x20000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="131"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x40000000" complete_cpuset="0x40000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="137" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x40000000" complete_cpuset="0x40000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="136" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="2" cpuset="0x40000000" complete_cpuset="0x40000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="134">
+              <object type="PU" os_index="30" cpuset="0x40000000" complete_cpuset="0x40000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="135"/>
+            </object>
+          </object>
+        </object>
+        <object type="L2Cache" cpuset="0x80000000" complete_cpuset="0x80000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="141" cache_size="524288" depth="2" cache_linesize="64" cache_associativity="16" cache_type="0">
+          <info name="Inclusive" value="0"/>
+          <object type="L1Cache" cpuset="0x80000000" complete_cpuset="0x80000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="140" cache_size="65536" depth="1" cache_linesize="64" cache_associativity="2" cache_type="1">
+            <info name="Inclusive" value="0"/>
+            <object type="Core" os_index="3" cpuset="0x80000000" complete_cpuset="0x80000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="138">
+              <object type="PU" os_index="31" cpuset="0x80000000" complete_cpuset="0x80000000" nodeset="0x0" complete_nodeset="0x00000080" gp_index="139"/>
+            </object>
+          </object>
+        </object>
+      </object>
+    </object>
+  </object>
+  <distances2 type="NUMANode" nbobjs="4" kind="5" indexing="os">
+    <indexes length="8">1 2 5 6 </indexes>
+    <u64values length="30">10 22 16 22 22 10 22 16 16 22 </u64values>
+    <u64values length="18">10 22 22 16 22 10 </u64values>
+  </distances2>
+</topology>
+

src/3rdparty/argon2/CMakeLists.txt

@@ -22,7 +22,7 @@ set(ARGON2_X86_64_SOURCES arch/x86_64/lib/argon2-arch.c arch/x86_64/lib/cpu-flag
 if (CMAKE_C_COMPILER_ID MATCHES MSVC)
     function(add_feature_impl FEATURE MSVC_FLAG DEF)
         add_library(argon2-${FEATURE} STATIC arch/x86_64/lib/argon2-${FEATURE}.c)
-        target_include_directories(argon2-${FEATURE} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include)
+        target_include_directories(argon2-${FEATURE} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../)
         target_include_directories(argon2-${FEATURE} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/lib)
         set_target_properties(argon2-${FEATURE} PROPERTIES POSITION_INDEPENDENT_CODE True)
 
@@ -38,7 +38,7 @@ if (CMAKE_C_COMPILER_ID MATCHES MSVC)
 elseif (NOT XMRIG_ARM AND CMAKE_SIZEOF_VOID_P EQUAL 8)
     function(add_feature_impl FEATURE GCC_FLAG DEF)
         add_library(argon2-${FEATURE} STATIC arch/x86_64/lib/argon2-${FEATURE}.c)
-        target_include_directories(argon2-${FEATURE} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include)
+        target_include_directories(argon2-${FEATURE} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../)
         target_include_directories(argon2-${FEATURE} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/lib)
         set_target_properties(argon2-${FEATURE} PROPERTIES POSITION_INDEPENDENT_CODE True)
 
@@ -84,5 +84,5 @@ endif()
 add_library(argon2 STATIC ${ARGON2_SOURCES})
 target_link_libraries(argon2 ${ARGON2_LIBS})
 
-target_include_directories(argon2 PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include)
+target_include_directories(argon2 PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../)
 target_include_directories(argon2 PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/lib)

src/3rdparty/argon2/lib/argon2.c

@@ -15,7 +15,7 @@
 #include <stdlib.h>
 #include <stdio.h>
 
-#include "argon2.h"
+#include "3rdparty/argon2.h"
 #include "encoding.h"
 #include "core.h"
 

src/3rdparty/argon2/lib/core.h

@@ -14,7 +14,7 @@
 #ifndef ARGON2_CORE_H
 #define ARGON2_CORE_H
 
-#include "argon2.h"
+#include "3rdparty/argon2.h"
 
 #if defined(_MSC_VER)
 #define ALIGN(n) __declspec(align(16))

src/3rdparty/argon2/lib/encoding.h

@@ -1,6 +1,6 @@
 #ifndef ENCODING_H
 #define ENCODING_H
-#include "argon2.h"
+#include "3rdparty/argon2.h"
 
 #define ARGON2_MAX_DECODED_LANES UINT32_C(255)
 #define ARGON2_MIN_DECODED_SALT_LEN UINT32_C(8)

src/3rdparty/argon2/lib/impl-select.c

@@ -3,7 +3,7 @@
 
 #include "impl-select.h"
 
-#include "argon2.h"
+#include "3rdparty/argon2.h"
 
 #define BENCH_SAMPLES 1024
 #define BENCH_MEM_BLOCKS 512

src/App.cpp

@@ -24,7 +24,7 @@
  */
 
 
-#include <stdlib.h>
+#include <cstdlib>
 #include <uv.h>
 
 
@@ -42,18 +42,9 @@
 #include "version.h"
 
 
-xmrig::App::App(Process *process) :
-    m_console(nullptr),
-    m_signals(nullptr)
+xmrig::App::App(Process *process)
 {
     m_controller = new Controller(process);
-    if (m_controller->init() != 0) {
-        return;
-    }
-
-    if (!m_controller->config()->isBackground()) {
-        m_console = new Console(this);
-    }
 }
 
 
@@ -68,12 +59,26 @@ xmrig::App::~App()
 int xmrig::App::exec()
 {
     if (!m_controller->isReady()) {
+        LOG_EMERG("no valid configuration found.");
+
         return 2;
     }
 
     m_signals = new Signals(this);
 
-    background();
+    int rc = 0;
+    if (background(rc)) {
+        return rc;
+    }
+
+    rc = m_controller->init();
+    if (rc != 0) {
+        return rc;
+    }
+
+    if (!m_controller->isBackground()) {
+        m_console = new Console(this);
+    }
 
     VirtualMemory::init(m_controller->config()->cpu().isHugePages());
 
@@ -87,10 +92,10 @@ int xmrig::App::exec()
 
     m_controller->start();
 
-    const int r = uv_run(uv_default_loop(), UV_RUN_DEFAULT);
+    rc = uv_run(uv_default_loop(), UV_RUN_DEFAULT);
     uv_loop_close(uv_default_loop());
 
-    return r;
+    return rc;
 }
 
 
@@ -150,7 +155,11 @@ void xmrig::App::onSignal(int signum)
 void xmrig::App::close()
 {
     m_signals->stop();
-    m_console->stop();
+
+    if (m_console) {
+        m_console->stop();
+    }
+
     m_controller->stop();
 
     Log::destroy();

src/App.h

@@ -29,6 +29,7 @@
 
 #include "base/kernel/interfaces/IConsoleListener.h"
 #include "base/kernel/interfaces/ISignalListener.h"
+#include "base/tools/Object.h"
 
 
 namespace xmrig {
@@ -44,6 +45,8 @@ class Signals;
 class App : public IConsoleListener, public ISignalListener
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(App)
+
     App(Process *process);
     ~App() override;
 
@@ -54,12 +57,12 @@ protected:
     void onSignal(int signum) override;
 
 private:
-    void background();
+    bool background(int &rc);
     void close();
 
-    Console *m_console;
-    Controller *m_controller;
-    Signals *m_signals;
+    Console *m_console          = nullptr;
+    Controller *m_controller    = nullptr;
+    Signals *m_signals          = nullptr;
 };
 
 

src/App_unix.cpp

@@ -23,33 +23,36 @@
  */
 
 
-#include <stdlib.h>
-#include <signal.h>
-#include <errno.h>
+#include <cstdlib>
+#include <csignal>
+#include <cerrno>
 #include <unistd.h>
 
 
 #include "App.h"
 #include "base/io/log/Log.h"
-#include "core/config/Config.h"
 #include "core/Controller.h"
 
 
-void xmrig::App::background()
+bool xmrig::App::background(int &rc)
 {
     signal(SIGPIPE, SIG_IGN);
 
-    if (!m_controller->config()->isBackground()) {
-        return;
+    if (!m_controller->isBackground()) {
+        return false;
     }
 
     int i = fork();
     if (i < 0) {
-        exit(1);
+        rc = 1;
+
+        return true;
     }
 
     if (i > 0) {
-        exit(0);
+        rc = 0;
+
+        return true;
     }
 
     i = setsid();
@@ -62,4 +65,6 @@ void xmrig::App::background()
     if (i < 0) {
         LOG_ERR("chdir() failed (errno = %d)", errno);
     }
+
+    return false;
 }

src/App_win.cpp

@@ -29,13 +29,12 @@
 
 #include "App.h"
 #include "core/Controller.h"
-#include "core/config/Config.h"
 
 
-void xmrig::App::background()
+bool xmrig::App::background(int &)
 {
-    if (!m_controller->config()->isBackground()) {
-        return;
+    if (!m_controller->isBackground()) {
+        return false;
     }
 
     HWND hcon = GetConsoleWindow();
@@ -46,4 +45,6 @@ void xmrig::App::background()
         CloseHandle(h);
         FreeConsole();
     }
+
+    return false;
 }

src/backend/common/Hashrate.cpp

@@ -23,10 +23,10 @@
  */
 
 
-#include <assert.h>
+#include <cassert>
 #include <cmath>
 #include <memory.h>
-#include <stdio.h>
+#include <cstdio>
 
 
 #include "backend/common/Hashrate.h"
@@ -133,8 +133,8 @@ double xmrig::Hashrate::calc(size_t threadId, size_t ms) const
         return nan("");
     }
 
-    const double hashes = static_cast<double>(lastestHashCnt - earliestHashCount);
-    const double time   = static_cast<double>(lastestStamp - earliestStamp) / 1000.0;
+    const auto hashes = static_cast<double>(lastestHashCnt - earliestHashCount);
+    const auto time   = static_cast<double>(lastestStamp - earliestStamp) / 1000.0;
 
     return hashes / time;
 }
@@ -175,3 +175,33 @@ rapidjson::Value xmrig::Hashrate::normalize(double d)
 
     return Value(floor(d * 100.0) / 100.0);
 }
+
+
+#ifdef XMRIG_FEATURE_API
+rapidjson::Value xmrig::Hashrate::toJSON(rapidjson::Document &doc) const
+{
+    using namespace rapidjson;
+    auto &allocator = doc.GetAllocator();
+
+    Value out(kArrayType);
+    out.PushBack(normalize(calc(ShortInterval)),  allocator);
+    out.PushBack(normalize(calc(MediumInterval)), allocator);
+    out.PushBack(normalize(calc(LargeInterval)),  allocator);
+
+    return out;
+}
+
+
+rapidjson::Value xmrig::Hashrate::toJSON(size_t threadId, rapidjson::Document &doc) const
+{
+    using namespace rapidjson;
+    auto &allocator = doc.GetAllocator();
+
+    Value out(kArrayType);
+    out.PushBack(normalize(calc(threadId, ShortInterval)),  allocator);
+    out.PushBack(normalize(calc(threadId, MediumInterval)), allocator);
+    out.PushBack(normalize(calc(threadId, LargeInterval)),  allocator);
+
+    return out;
+}
+#endif

src/backend/common/Hashrate.h

@@ -26,10 +26,11 @@
 #define XMRIG_HASHRATE_H
 
 
-#include <stddef.h>
-#include <stdint.h>
+#include <cstddef>
+#include <cstdint>
 
 
+#include "base/tools/Object.h"
 #include "rapidjson/fwd.h"
 
 
@@ -39,6 +40,8 @@ namespace xmrig {
 class Hashrate
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(Hashrate)
+
     enum Intervals {
         ShortInterval  = 10000,
         MediumInterval = 60000,
@@ -58,6 +61,11 @@ public:
     static const char *format(double h, char *buf, size_t size);
     static rapidjson::Value normalize(double d);
 
+#   ifdef XMRIG_FEATURE_API
+    rapidjson::Value toJSON(rapidjson::Document &doc) const;
+    rapidjson::Value toJSON(size_t threadId, rapidjson::Document &doc) const;
+#   endif
+
 private:
     constexpr static size_t kBucketSize = 2 << 11;
     constexpr static size_t kBucketMask = kBucketSize - 1;

src/backend/common/Thread.h

@@ -26,10 +26,11 @@
 #define XMRIG_THREAD_H
 
 
-#include <thread>
-
-
 #include "backend/common/interfaces/IWorker.h"
+#include "base/tools/Object.h"
+
+
+#include <thread>
 
 
 namespace xmrig {
@@ -42,6 +43,8 @@ template<class T>
 class Thread
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(Thread)
+
     inline Thread(IBackend *backend, size_t id, const T &config) : m_id(id), m_config(config), m_backend(backend) {}
     inline ~Thread() { m_thread.join(); delete m_worker; }
 

src/backend/common/Threads.h

@@ -44,6 +44,7 @@ class Threads
 public:
     inline bool has(const char *profile) const                                         { return m_profiles.count(profile) > 0; }
     inline bool isDisabled(const Algorithm &algo) const                                { return m_disabled.count(algo) > 0; }
+    inline bool isEmpty() const                                                        { return m_profiles.empty(); }
     inline bool isExist(const Algorithm &algo) const                                   { return isDisabled(algo) || m_aliases.count(algo) > 0 || has(algo.shortName()); }
     inline const T &get(const Algorithm &algo, bool strict = false) const              { return get(profileName(algo, strict)); }
     inline void disable(const Algorithm &algo)                                         { m_disabled.insert(algo); }

src/backend/common/Workers.cpp

@@ -29,6 +29,7 @@
 #include "backend/common/Workers.h"
 #include "backend/cpu/CpuWorker.h"
 #include "base/io/log/Log.h"
+#include "base/tools/Object.h"
 
 
 #ifdef XMRIG_FEATURE_OPENCL
@@ -42,9 +43,10 @@ namespace xmrig {
 class WorkersPrivate
 {
 public:
-    inline WorkersPrivate()
-    {
-    }
+    XMRIG_DISABLE_COPY_MOVE(WorkersPrivate)
+
+
+    WorkersPrivate() = default;
 
 
     inline ~WorkersPrivate()
@@ -98,6 +100,7 @@ void xmrig::Workers<T>::start(const std::vector<T> &data)
     }
 
     d_ptr->hashrate = new Hashrate(m_workers.size());
+    Nonce::touch(T::backend());
 
     for (Thread<T> *worker : m_workers) {
         worker->start(Workers<T>::onReady);
@@ -131,7 +134,7 @@ void xmrig::Workers<T>::tick(uint64_t)
 
     for (Thread<T> *handle : m_workers) {
         if (!handle->worker()) {
-            return;
+            continue;
         }
 
         d_ptr->hashrate->add(handle->id(), handle->worker()->hashCount(), handle->worker()->timestamp());
@@ -154,17 +157,21 @@ void xmrig::Workers<T>::onReady(void *arg)
     auto handle = static_cast<Thread<T>* >(arg);
 
     IWorker *worker = create(handle);
-    if (!worker || !worker->selfTest()) {
-        LOG_ERR("thread %zu error: \"hash self-test failed\".", worker->id());
+    assert(worker != nullptr);
 
+    if (!worker || !worker->selfTest()) {
+        LOG_ERR("%s " RED("thread ") RED_BOLD("#%zu") RED(" self-test failed"), T::tag(), worker->id());
+
+        handle->backend()->start(worker, false);
         delete worker;
+
         return;
     }
 
     assert(handle->backend() != nullptr);
 
     handle->setWorker(worker);
-    handle->backend()->start(worker);
+    handle->backend()->start(worker, true);
 }
 
 

src/backend/common/Workers.h

@@ -29,6 +29,7 @@
 
 #include "backend/common/Thread.h"
 #include "backend/cpu/CpuLaunchData.h"
+#include "base/tools/Object.h"
 
 
 #ifdef XMRIG_FEATURE_OPENCL
@@ -47,6 +48,8 @@ template<class T>
 class Workers
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE(Workers)
+
     Workers();
     ~Workers();
 

src/backend/common/interfaces/IBackend.h

@@ -26,7 +26,7 @@
 #define XMRIG_IBACKEND_H
 
 
-#include <stdint.h>
+#include <cstdint>
 
 
 #include "rapidjson/fwd.h"
@@ -56,7 +56,7 @@ public:
     virtual void prepare(const Job &nextJob)                            = 0;
     virtual void printHashrate(bool details)                            = 0;
     virtual void setJob(const Job &job)                                 = 0;
-    virtual void start(IWorker *worker)                                 = 0;
+    virtual void start(IWorker *worker, bool ready)                     = 0;
     virtual void stop()                                                 = 0;
     virtual void tick(uint64_t ticks)                                   = 0;
 

src/backend/common/interfaces/IWorker.h

@@ -26,8 +26,8 @@
 #define XMRIG_IWORKER_H
 
 
-#include <stdint.h>
-#include <stddef.h>
+#include <cstdint>
+#include <cstddef>
 
 
 namespace xmrig {
@@ -44,6 +44,7 @@ public:
     virtual bool selfTest()                         = 0;
     virtual const VirtualMemory *memory() const     = 0;
     virtual size_t id() const                       = 0;
+    virtual size_t intensity() const                = 0;
     virtual uint64_t hashCount() const              = 0;
     virtual uint64_t timestamp() const              = 0;
     virtual void start()                            = 0;

src/backend/common/misc/PciTopology.h

@@ -27,7 +27,7 @@
 #define XMRIG_PCITOPOLOGY_H
 
 
-#include <stdio.h>
+#include <cstdio>
 
 
 #include "base/tools/String.h"
@@ -40,19 +40,30 @@ class PciTopology
 {
 public:
     PciTopology() = default;
-    PciTopology(uint32_t bus, uint32_t device, uint32_t function) : bus(bus), device(device), function(function) {}
+    PciTopology(uint32_t bus, uint32_t device, uint32_t function) : m_valid(true), m_bus(bus), m_device(device), m_function(function) {}
 
-    uint32_t bus        = 0;
-    uint32_t device     = 0;
-    uint32_t function   = 0;
+    inline bool isValid() const        { return m_valid; }
+    inline uint8_t bus() const         { return m_bus; }
+    inline uint8_t device() const      { return m_device; }
+    inline uint8_t function() const    { return m_function; }
 
     String toString() const
     {
+        if (!isValid()) {
+            return "n/a";
+        }
+
         char *buf = new char[8]();
-        snprintf(buf, 8, "%02x:%02x.%01x", bus, device, function);
+        snprintf(buf, 8, "%02hhx:%02hhx.%01hhx", bus(), device(), function());
 
         return buf;
     }
+
+private:
+    bool m_valid         = false;
+    uint8_t m_bus        = 0;
+    uint8_t m_device     = 0;
+    uint8_t m_function   = 0;
 };
 
 

src/backend/cpu/Cpu.cpp

@@ -23,7 +23,7 @@
  */
 
 
-#include <assert.h>
+#include <cassert>
 
 
 #include "backend/cpu/Cpu.h"
@@ -44,7 +44,15 @@ static xmrig::ICpuInfo *cpuInfo = nullptr;
 
 xmrig::ICpuInfo *xmrig::Cpu::info()
 {
-    assert(cpuInfo != nullptr);
+    if (cpuInfo == nullptr) {
+#       if defined(XMRIG_FEATURE_HWLOC)
+        cpuInfo = new HwlocCpuInfo();
+#       elif defined(XMRIG_FEATURE_LIBCPUID)
+        cpuInfo = new AdvancedCpuInfo();
+#       else
+        cpuInfo = new BasicCpuInfo();
+#       endif
+    }
 
     return cpuInfo;
 }
@@ -62,7 +70,7 @@ rapidjson::Value xmrig::Cpu::toJSON(rapidjson::Document &doc)
     cpu.AddMember("brand",      StringRef(i->brand()), allocator);
     cpu.AddMember("aes",        i->hasAES(), allocator);
     cpu.AddMember("avx2",       i->hasAVX2(), allocator);
-    cpu.AddMember("x64",        i->isX64(), allocator);
+    cpu.AddMember("x64",        ICpuInfo::isX64(), allocator);
     cpu.AddMember("l2",         static_cast<uint64_t>(i->L2()), allocator);
     cpu.AddMember("l3",         static_cast<uint64_t>(i->L3()), allocator);
     cpu.AddMember("cores",      static_cast<uint64_t>(i->cores()), allocator);
@@ -81,20 +89,6 @@ rapidjson::Value xmrig::Cpu::toJSON(rapidjson::Document &doc)
 }
 
 
-void xmrig::Cpu::init()
-{
-    assert(cpuInfo == nullptr);
-
-#   if defined(XMRIG_FEATURE_HWLOC)
-    cpuInfo = new HwlocCpuInfo();
-#   elif defined(XMRIG_FEATURE_LIBCPUID)
-    cpuInfo = new AdvancedCpuInfo();
-#   else
-    cpuInfo = new BasicCpuInfo();
-#   endif
-}
-
-
 void xmrig::Cpu::release()
 {
     assert(cpuInfo != nullptr);

src/backend/cpu/Cpu.h

@@ -37,7 +37,6 @@ class Cpu
 public:
     static ICpuInfo *info();
     static rapidjson::Value toJSON(rapidjson::Document &doc);
-    static void init();
     static void release();
 
     inline static Assembly::Id assembly(Assembly::Id hint) { return hint == Assembly::AUTO ? Cpu::info()->assembly() : hint; }

src/backend/cpu/CpuBackend.cpp

@@ -60,7 +60,7 @@ namespace xmrig {
 extern template class Threads<CpuThreads>;
 
 
-static const char *tag      = CYAN_BG_BOLD(" cpu ");
+static const char *tag      = CYAN_BG_BOLD(WHITE_BOLD_S " cpu ");
 static const String kType   = "cpu";
 static std::mutex mutex;
 
@@ -80,38 +80,51 @@ public:
         m_memory    = memory;
         m_pages     = 0;
         m_started   = 0;
+        m_errors    = 0;
         m_threads   = threads.size();
         m_ways      = 0;
         m_ts        = Chrono::steadyMSecs();
-
-        for (const CpuLaunchData &data : threads) {
-            m_ways += data.intensity;
-        }
     }
 
-    inline bool started(const std::pair<size_t, size_t> &hugePages)
+    inline bool started(IWorker *worker, bool ready)
     {
-        m_started++;
-        m_hugePages += hugePages.first;
-        m_pages     += hugePages.second;
+        if (ready) {
+            auto hugePages = worker->memory()->hugePages();
 
-        return m_started == m_threads;
+            m_started++;
+            m_hugePages += hugePages.first;
+            m_pages     += hugePages.second;
+            m_ways      += worker->intensity();
+        }
+        else {
+            m_errors++;
+        }
+
+        return (m_started + m_errors) == m_threads;
     }
 
     inline void print() const
     {
-        LOG_INFO("%s" GREEN_BOLD(" READY") " threads " CYAN_BOLD("%zu(%zu)") " huge pages %s%zu/%zu %1.0f%%\x1B[0m memory " CYAN_BOLD("%zu KB") BLACK_BOLD(" (%" PRIu64 " ms)"),
+        if (m_started == 0) {
+            LOG_ERR("%s " RED_BOLD("disabled") YELLOW(" (failed to start threads)"), tag);
+
+            return;
+        }
+
+        LOG_INFO("%s" GREEN_BOLD(" READY") " threads %s%zu/%zu (%zu)" CLEAR " huge pages %s%zu/%zu %1.0f%%" CLEAR " memory " CYAN_BOLD("%zu KB") BLACK_BOLD(" (%" PRIu64 " ms)"),
                  tag,
-                 m_threads, m_ways,
+                 m_errors == 0 ? CYAN_BOLD_S : YELLOW_BOLD_S,
+                 m_started, m_threads, m_ways,
                  (m_hugePages == m_pages ? GREEN_BOLD_S : (m_hugePages == 0 ? RED_BOLD_S : YELLOW_BOLD_S)),
                  m_hugePages, m_pages,
                  m_hugePages == 0 ? 0.0 : static_cast<double>(m_hugePages) / m_pages * 100.0,
-                 m_ways * m_memory / 1024,
+                 memory() / 1024,
                  Chrono::steadyMSecs() - m_ts
                  );
     }
 
 private:
+    size_t m_errors       = 0;
     size_t m_hugePages    = 0;
     size_t m_memory       = 0;
     size_t m_pages        = 0;
@@ -322,17 +335,19 @@ void xmrig::CpuBackend::setJob(const Job &job)
 }
 
 
-void xmrig::CpuBackend::start(IWorker *worker)
+void xmrig::CpuBackend::start(IWorker *worker, bool ready)
 {
     mutex.lock();
 
-    if (d_ptr->status.started(worker->memory()->hugePages())) {
+    if (d_ptr->status.started(worker, ready)) {
         d_ptr->status.print();
     }
 
     mutex.unlock();
 
-    worker->start();
+    if (ready) {
+        worker->start();
+    }
 }
 
 
@@ -390,8 +405,9 @@ rapidjson::Value xmrig::CpuBackend::toJSON(rapidjson::Document &doc) const
         return out;
     }
 
+    out.AddMember("hashrate", hashrate()->toJSON(doc), allocator);
+
     Value threads(kArrayType);
-    const Hashrate *hr = hashrate();
 
     size_t i = 0;
     for (const CpuLaunchData &data : d_ptr->threads) {
@@ -399,15 +415,9 @@ rapidjson::Value xmrig::CpuBackend::toJSON(rapidjson::Document &doc) const
         thread.AddMember("intensity",   data.intensity, allocator);
         thread.AddMember("affinity",    data.affinity, allocator);
         thread.AddMember("av",          data.av(), allocator);
-
-        Value hashrate(kArrayType);
-        hashrate.PushBack(Hashrate::normalize(hr->calc(i, Hashrate::ShortInterval)),  allocator);
-        hashrate.PushBack(Hashrate::normalize(hr->calc(i, Hashrate::MediumInterval)), allocator);
-        hashrate.PushBack(Hashrate::normalize(hr->calc(i, Hashrate::LargeInterval)),  allocator);
+        thread.AddMember("hashrate",    hashrate()->toJSON(i, doc), allocator);
 
         i++;
-
-        thread.AddMember("hashrate", hashrate, allocator);
         threads.PushBack(thread, allocator);
     }
 

src/backend/cpu/CpuBackend.h

@@ -26,10 +26,11 @@
 #define XMRIG_CPUBACKEND_H
 
 
-#include <utility>
-
-
 #include "backend/common/interfaces/IBackend.h"
+#include "base/tools/Object.h"
+
+
+#include <utility>
 
 
 namespace xmrig {
@@ -43,6 +44,8 @@ class Miner;
 class CpuBackend : public IBackend
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(CpuBackend)
+
     CpuBackend(Controller *controller);
     ~CpuBackend() override;
 
@@ -55,7 +58,7 @@ protected:
     void prepare(const Job &nextJob) override;
     void printHashrate(bool details) override;
     void setJob(const Job &job) override;
-    void start(IWorker *worker) override;
+    void start(IWorker *worker, bool ready) override;
     void stop() override;
     void tick(uint64_t ticks) override;
 

src/backend/cpu/CpuConfig.cpp

@@ -35,6 +35,7 @@ static const char *kCn                  = "cn";
 static const char *kEnabled             = "enabled";
 static const char *kHugePages           = "huge-pages";
 static const char *kHwAes               = "hw-aes";
+static const char *kMaxThreadsHint      = "max-threads-hint";
 static const char *kPriority            = "priority";
 
 #ifdef XMRIG_FEATURE_ASM
@@ -72,11 +73,6 @@ extern template class Threads<CpuThreads>;
 }
 
 
-xmrig::CpuConfig::CpuConfig()
-{
-}
-
-
 bool xmrig::CpuConfig::isHwAES() const
 {
     return (m_aes == AES_AUTO ? (Cpu::info()->hasAES() ? AES_HW : AES_SOFT) : m_aes) == AES_HW;
@@ -95,6 +91,10 @@ rapidjson::Value xmrig::CpuConfig::toJSON(rapidjson::Document &doc) const
     obj.AddMember(StringRef(kHwAes),        m_aes == AES_AUTO ? Value(kNullType) : Value(m_aes == AES_HW), allocator);
     obj.AddMember(StringRef(kPriority),     priority() != -1 ? Value(priority()) : Value(kNullType), allocator);
 
+    if (m_threads.isEmpty()) {
+        obj.AddMember(StringRef(kMaxThreadsHint), m_limit, allocator);
+    }
+
 #   ifdef XMRIG_FEATURE_ASM
     obj.AddMember(StringRef(kAsm), m_assembly.toJSON(), allocator);
 #   endif
@@ -131,8 +131,9 @@ std::vector<xmrig::CpuLaunchData> xmrig::CpuConfig::get(const Miner *miner, cons
 void xmrig::CpuConfig::read(const rapidjson::Value &value, uint32_t version)
 {
     if (value.IsObject()) {
-        m_enabled       = Json::getBool(value, kEnabled, m_enabled);
-        m_hugePages     = Json::getBool(value, kHugePages, m_hugePages);
+        m_enabled   = Json::getBool(value, kEnabled, m_enabled);
+        m_hugePages = Json::getBool(value, kHugePages, m_hugePages);
+        m_limit     = Json::getUint(value, kMaxThreadsHint, m_limit);
 
         setAesMode(Json::getValue(value, kHwAes));
         setPriority(Json::getInt(value,  kPriority, -1));
@@ -168,28 +169,28 @@ void xmrig::CpuConfig::generate()
     ICpuInfo *cpu = Cpu::info();
 
     m_threads.disable(Algorithm::CN_0);
-    m_threads.move(kCn, cpu->threads(Algorithm::CN_0));
+    m_threads.move(kCn, cpu->threads(Algorithm::CN_0, m_limit));
 
 #   ifdef XMRIG_ALGO_CN_GPU
-    m_threads.move(kCnGPU, cpu->threads(Algorithm::CN_GPU));
+    m_threads.move(kCnGPU, cpu->threads(Algorithm::CN_GPU, m_limit));
 #   endif
 
 #   ifdef XMRIG_ALGO_CN_LITE
     m_threads.disable(Algorithm::CN_LITE_0);
-    m_threads.move(kCnLite, cpu->threads(Algorithm::CN_LITE_1));
+    m_threads.move(kCnLite, cpu->threads(Algorithm::CN_LITE_1, m_limit));
 #   endif
 
 #   ifdef XMRIG_ALGO_CN_HEAVY
-    m_threads.move(kCnHeavy, cpu->threads(Algorithm::CN_HEAVY_0));
+    m_threads.move(kCnHeavy, cpu->threads(Algorithm::CN_HEAVY_0, m_limit));
 #   endif
 
 #   ifdef XMRIG_ALGO_CN_PICO
-    m_threads.move(kCnPico, cpu->threads(Algorithm::CN_PICO_0));
+    m_threads.move(kCnPico, cpu->threads(Algorithm::CN_PICO_0, m_limit));
 #   endif
 
 #   ifdef XMRIG_ALGO_RANDOMX
-    m_threads.move(kRx, cpu->threads(Algorithm::RX_0));
-    m_threads.move(kRxWOW, cpu->threads(Algorithm::RX_WOW));
+    m_threads.move(kRx, cpu->threads(Algorithm::RX_0, m_limit));
+    m_threads.move(kRxWOW, cpu->threads(Algorithm::RX_WOW, m_limit));
 #   endif
 
     generateArgon2();
@@ -199,7 +200,7 @@ void xmrig::CpuConfig::generate()
 void xmrig::CpuConfig::generateArgon2()
 {
 #   ifdef XMRIG_ALGO_ARGON2
-    m_threads.move(kArgon2, Cpu::info()->threads(Algorithm::AR2_CHUKWA));
+    m_threads.move(kArgon2, Cpu::info()->threads(Algorithm::AR2_CHUKWA, m_limit));
 #   endif
 }
 

src/backend/cpu/CpuConfig.h

@@ -44,7 +44,7 @@ public:
         AES_SOFT
     };
 
-    CpuConfig();
+    CpuConfig() = default;
 
     bool isHwAES() const;
     rapidjson::Value toJSON(rapidjson::Document &doc) const;
@@ -74,6 +74,7 @@ private:
     int m_priority       = -1;
     String m_argon2Impl;
     Threads<CpuThreads> m_threads;
+    uint32_t m_limit     = 100;
 };
 
 

src/backend/cpu/CpuLaunchData.cpp

@@ -24,13 +24,15 @@
  */
 
 
-#include <algorithm>
-
-
 #include "backend/cpu/CpuLaunchData.h"
+
+#include "backend/common/Tags.h"
 #include "backend/cpu/CpuConfig.h"
 
 
+#include <algorithm>
+
+
 xmrig::CpuLaunchData::CpuLaunchData(const Miner *miner, const Algorithm &algorithm, const CpuConfig &config, const CpuThread &thread) :
     algorithm(algorithm),
     assembly(config.assembly()),
@@ -65,3 +67,9 @@ xmrig::CnHash::AlgoVariant xmrig::CpuLaunchData::av() const
 
     return static_cast<CnHash::AlgoVariant>(!hwAES ? (intensity + 5) : (intensity + 2));
 }
+
+
+const char *xmrig::CpuLaunchData::tag()
+{
+    return cpu_tag();
+}

src/backend/cpu/CpuLaunchData.h

@@ -54,6 +54,8 @@ public:
     inline bool operator!=(const CpuLaunchData &other) const    { return !isEqual(other); }
     inline bool operator==(const CpuLaunchData &other) const    { return isEqual(other); }
 
+    static const char *tag();
+
     const Algorithm algorithm;
     const Assembly assembly;
     const bool hugePages;

src/backend/cpu/CpuWorker.h

@@ -54,6 +54,7 @@ protected:
     void start() override;
 
     inline const VirtualMemory *memory() const override { return m_memory; }
+    inline size_t intensity() const override            { return N; }
 
 private:
     inline cn_hash_fun fn(const Algorithm &algorithm) const { return CnHash::fn(algorithm, m_av, m_assembly); }

src/backend/cpu/interfaces/ICpuInfo.h

@@ -45,18 +45,18 @@ public:
     inline constexpr static bool isX64() { return false; }
 #   endif
 
-    virtual Assembly::Id assembly() const                                     = 0;
-    virtual bool hasAES() const                                               = 0;
-    virtual bool hasAVX2() const                                              = 0;
-    virtual const char *backend() const                                       = 0;
-    virtual const char *brand() const                                         = 0;
-    virtual CpuThreads threads(const Algorithm &algorithm) const              = 0;
-    virtual size_t cores() const                                              = 0;
-    virtual size_t L2() const                                                 = 0;
-    virtual size_t L3() const                                                 = 0;
-    virtual size_t nodes() const                                              = 0;
-    virtual size_t packages() const                                           = 0;
-    virtual size_t threads() const                                            = 0;
+    virtual Assembly::Id assembly() const                                           = 0;
+    virtual bool hasAES() const                                                     = 0;
+    virtual bool hasAVX2() const                                                    = 0;
+    virtual const char *backend() const                                             = 0;
+    virtual const char *brand() const                                               = 0;
+    virtual CpuThreads threads(const Algorithm &algorithm, uint32_t limit) const    = 0;
+    virtual size_t cores() const                                                    = 0;
+    virtual size_t L2() const                                                       = 0;
+    virtual size_t L3() const                                                       = 0;
+    virtual size_t nodes() const                                                    = 0;
+    virtual size_t packages() const                                                 = 0;
+    virtual size_t threads() const                                                  = 0;
 };
 
 

src/backend/cpu/platform/AdvancedCpuInfo.cpp

@@ -23,10 +23,10 @@
  */
 
 #include <algorithm>
-#include <assert.h>
-#include <math.h>
-#include <stdio.h>
-#include <string.h>
+#include <cassert>
+#include <cmath>
+#include <cstdio>
+#include <cstring>
 
 
 #include "3rdparty/libcpuid/libcpuid.h"
@@ -109,7 +109,7 @@ xmrig::AdvancedCpuInfo::AdvancedCpuInfo() :
 }
 
 
-xmrig::CpuThreads xmrig::AdvancedCpuInfo::threads(const Algorithm &algorithm) const
+xmrig::CpuThreads xmrig::AdvancedCpuInfo::threads(const Algorithm &algorithm, uint32_t limit) const
 {
     if (threads() == 1) {
         return 1;
@@ -153,5 +153,12 @@ xmrig::CpuThreads xmrig::AdvancedCpuInfo::threads(const Algorithm &algorithm) co
     }
 #   endif
 
-    return CpuThreads(std::max<size_t>(std::min<size_t>(count, threads()), 1), intensity);
+    if (limit > 0 && limit < 100) {
+        count = std::min(count, static_cast<size_t>(round(threads() * (limit / 100.0))));
+    }
+    else {
+        count = std::min(count, threads());
+    }
+
+    return CpuThreads(std::max<size_t>(count, 1), intensity);
 }

src/backend/cpu/platform/AdvancedCpuInfo.h

@@ -38,7 +38,7 @@ public:
     AdvancedCpuInfo();
 
 protected:
-    CpuThreads threads(const Algorithm &algorithm) const override;
+    CpuThreads threads(const Algorithm &algorithm, uint32_t limit) const override;
 
     inline Assembly::Id assembly() const override   { return m_assembly; }
     inline bool hasAES() const override             { return m_aes; }

src/backend/cpu/platform/BasicCpuInfo.cpp

@@ -179,7 +179,7 @@ const char *xmrig::BasicCpuInfo::backend() const
 }
 
 
-xmrig::CpuThreads xmrig::BasicCpuInfo::threads(const Algorithm &algorithm) const
+xmrig::CpuThreads xmrig::BasicCpuInfo::threads(const Algorithm &algorithm, uint32_t limit) const
 {
     const size_t count = std::thread::hardware_concurrency();
 

src/backend/cpu/platform/BasicCpuInfo.h

@@ -39,7 +39,7 @@ public:
 
 protected:
     const char *backend() const override;
-    CpuThreads threads(const Algorithm &algorithm) const override;
+    CpuThreads threads(const Algorithm &algorithm, uint32_t limit) const override;
 
     inline Assembly::Id assembly() const override   { return m_assembly; }
     inline bool hasAES() const override             { return m_aes; }

src/backend/cpu/platform/BasicCpuInfo_arm.cpp

@@ -63,7 +63,7 @@ const char *xmrig::BasicCpuInfo::backend() const
 }
 
 
-xmrig::CpuThreads xmrig::BasicCpuInfo::threads(const Algorithm &) const
+xmrig::CpuThreads xmrig::BasicCpuInfo::threads(const Algorithm &, uint32_t) const
 {
     return CpuThreads(threads());
 }

src/backend/cpu/platform/HwlocCpuInfo.cpp

@@ -29,6 +29,7 @@
 
 
 #include <algorithm>
+#include <cmath>
 #include <hwloc.h>
 
 
@@ -127,9 +128,7 @@ static inline bool isCacheExclusive(hwloc_obj_t obj)
 } // namespace xmrig
 
 
-xmrig::HwlocCpuInfo::HwlocCpuInfo() : BasicCpuInfo(),
-    m_backend(),
-    m_cache()
+xmrig::HwlocCpuInfo::HwlocCpuInfo()
 {
     m_threads = 0;
 
@@ -149,7 +148,7 @@ xmrig::HwlocCpuInfo::HwlocCpuInfo() : BasicCpuInfo(),
 #   endif
 
     const std::vector<hwloc_obj_t> packages = findByType(hwloc_get_root_obj(m_topology), HWLOC_OBJ_PACKAGE);
-    if (packages.size()) {
+    if (!packages.empty()) {
         const char *value = hwloc_obj_get_info_by_name(packages[0], "CPUModel");
         if (value) {
             strncpy(m_brand, value, 64);
@@ -178,7 +177,7 @@ xmrig::HwlocCpuInfo::HwlocCpuInfo() : BasicCpuInfo(),
 
     m_threads   = countByType(m_topology, HWLOC_OBJ_PU);
     m_cores     = countByType(m_topology, HWLOC_OBJ_CORE);
-    m_nodes     = std::max<size_t>(countByType(m_topology, HWLOC_OBJ_NUMANODE), 1);
+    m_nodes     = std::max(hwloc_bitmap_weight(hwloc_topology_get_complete_nodeset(m_topology)), 1);
     m_packages  = countByType(m_topology, HWLOC_OBJ_PACKAGE);
 
     if (m_nodes > 1) {
@@ -202,10 +201,10 @@ xmrig::HwlocCpuInfo::~HwlocCpuInfo()
 }
 
 
-xmrig::CpuThreads xmrig::HwlocCpuInfo::threads(const Algorithm &algorithm) const
+xmrig::CpuThreads xmrig::HwlocCpuInfo::threads(const Algorithm &algorithm, uint32_t limit) const
 {
     if (L2() == 0 && L3() == 0) {
-        return BasicCpuInfo::threads(algorithm);
+        return BasicCpuInfo::threads(algorithm, limit);
     }
 
     const unsigned depth = L3() > 0 ? 3 : 2;
@@ -218,21 +217,37 @@ xmrig::CpuThreads xmrig::HwlocCpuInfo::threads(const Algorithm &algorithm) const
 
     findCache(hwloc_get_root_obj(m_topology), depth, depth, [&caches](hwloc_obj_t found) { caches.emplace_back(found); });
 
-    for (hwloc_obj_t cache : caches) {
-        processTopLevelCache(cache, algorithm, threads);
+    if (limit > 0 && limit < 100 && !caches.empty()) {
+        const double maxTotalThreads = round(m_threads * (limit / 100.0));
+        const auto maxPerCache       = std::max(static_cast<int>(round(maxTotalThreads / caches.size())), 1);
+        int remaining                = std::max(static_cast<int>(maxTotalThreads), 1);
+
+        for (hwloc_obj_t cache : caches) {
+            processTopLevelCache(cache, algorithm, threads, std::min(maxPerCache, remaining));
+
+            remaining -= maxPerCache;
+            if (remaining <= 0) {
+                break;
+            }
+        }
+    }
+    else {
+        for (hwloc_obj_t cache : caches) {
+            processTopLevelCache(cache, algorithm, threads, 0);
+        }
     }
 
     if (threads.isEmpty()) {
         LOG_WARN("hwloc auto configuration for algorithm \"%s\" failed.", algorithm.shortName());
 
-        return BasicCpuInfo::threads(algorithm);
+        return BasicCpuInfo::threads(algorithm, limit);
     }
 
     return threads;
 }
 
 
-void xmrig::HwlocCpuInfo::processTopLevelCache(hwloc_obj_t cache, const Algorithm &algorithm, CpuThreads &threads) const
+void xmrig::HwlocCpuInfo::processTopLevelCache(hwloc_obj_t cache, const Algorithm &algorithm, CpuThreads &threads, size_t limit) const
 {
     constexpr size_t oneMiB = 1024u * 1024u;
 
@@ -296,6 +311,10 @@ void xmrig::HwlocCpuInfo::processTopLevelCache(hwloc_obj_t cache, const Algorith
     }
 #   endif
 
+    if (limit > 0) {
+        cacheHashes = std::min(cacheHashes, limit);
+    }
+
     if (cacheHashes >= PUs) {
         for (hwloc_obj_t core : cores) {
             const std::vector<hwloc_obj_t> units = findByType(core, HWLOC_OBJ_PU);

src/backend/cpu/platform/HwlocCpuInfo.h

@@ -27,10 +27,11 @@
 
 
 #include "backend/cpu/platform/BasicCpuInfo.h"
+#include "base/tools/Object.h"
 
 
-typedef struct hwloc_obj *hwloc_obj_t;
-typedef struct hwloc_topology *hwloc_topology_t;
+using hwloc_obj_t       = struct hwloc_obj *;
+using hwloc_topology_t  = struct hwloc_topology *;
 
 
 namespace xmrig {
@@ -39,6 +40,9 @@ namespace xmrig {
 class HwlocCpuInfo : public BasicCpuInfo
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE(HwlocCpuInfo)
+
+
     enum Feature : uint32_t {
         SET_THISTHREAD_MEMBIND = 1
     };
@@ -51,7 +55,7 @@ public:
     static inline const std::vector<uint32_t> &nodeIndexes()    { return m_nodeIndexes; }
 
 protected:
-    CpuThreads threads(const Algorithm &algorithm) const override;
+    CpuThreads threads(const Algorithm &algorithm, uint32_t limit) const override;
 
     inline const char *backend() const override     { return m_backend; }
     inline size_t cores() const override            { return m_cores; }
@@ -61,17 +65,17 @@ protected:
     inline size_t packages() const override         { return m_packages; }
 
 private:
-    void processTopLevelCache(hwloc_obj_t obj, const Algorithm &algorithm, CpuThreads &threads) const;
+    void processTopLevelCache(hwloc_obj_t obj, const Algorithm &algorithm, CpuThreads &threads, size_t limit) const;
 
     static std::vector<uint32_t> m_nodeIndexes;
     static uint32_t m_features;
 
-    char m_backend[20];
-    hwloc_topology_t m_topology;
-    size_t m_cache[5];
-    size_t m_cores      = 0;
-    size_t m_nodes      = 0;
-    size_t m_packages   = 0;
+    char m_backend[20]          = { 0 };
+    hwloc_topology_t m_topology = nullptr;
+    size_t m_cache[5]           = { 0 };
+    size_t m_cores              = 0;
+    size_t m_nodes              = 0;
+    size_t m_packages           = 0;
 };
 
 

src/backend/opencl/OclBackend.cpp

@@ -35,6 +35,7 @@
 #include "backend/opencl/OclConfig.h"
 #include "backend/opencl/OclLaunchData.h"
 #include "backend/opencl/OclWorker.h"
+#include "backend/opencl/runners/tools/OclSharedState.h"
 #include "backend/opencl/wrappers/OclContext.h"
 #include "backend/opencl/wrappers/OclLib.h"
 #include "base/io/log/Log.h"
@@ -72,12 +73,19 @@ static void printDisabled(const char *reason)
 struct OclLaunchStatus
 {
 public:
-    inline bool started()               { m_started++; return m_started == m_threads; }
-    inline size_t threads() const       { return m_threads; }
+    inline size_t threads() const { return m_threads; }
+
+    inline bool started(bool ready)
+    {
+        ready ? m_started++ : m_errors++;
+
+        return (m_started + m_errors) == m_threads;
+    }
 
     inline void start(size_t threads)
     {
         m_started        = 0;
+        m_errors         = 0;
         m_threads        = threads;
         m_ts             = Chrono::steadyMSecs();
         OclWorker::ready = false;
@@ -85,14 +93,23 @@ public:
 
     inline void print() const
     {
-        LOG_INFO("%s" GREEN_BOLD(" READY") " threads " CYAN_BOLD("%zu") BLACK_BOLD(" (%" PRIu64 " ms)"),
+        if (m_started == 0) {
+            LOG_ERR("%s " RED_BOLD("disabled") YELLOW(" (failed to start threads)"), tag);
+
+            return;
+        }
+
+        LOG_INFO("%s" GREEN_BOLD(" READY") " threads " "%s%zu/%zu" BLACK_BOLD(" (%" PRIu64 " ms)"),
                  tag,
+                 m_errors == 0 ? CYAN_BOLD_S : YELLOW_BOLD_S,
+                 m_started,
                  m_threads,
                  Chrono::steadyMSecs() - m_ts
                  );
     }
 
 private:
+    size_t m_errors     = 0;
     size_t m_started    = 0;
     size_t m_threads    = 0;
     uint64_t m_ts       = 0;
@@ -119,6 +136,10 @@ public:
             return printDisabled(RED_S " (failed to load OpenCL runtime)");
         }
 
+        if (platform.isValid()) {
+            return;
+        }
+
         platform = cl.platform();
         if (!platform.isValid()) {
             return printDisabled(RED_S " (selected OpenCL platform NOT found)");
@@ -134,7 +155,7 @@ public:
         for (const OclDevice &device : devices) {
             Log::print(GREEN_BOLD(" * ") WHITE_BOLD("%-13s") CYAN_BOLD("#%zu") YELLOW(" %s") " %s " WHITE_BOLD("%uMHz") " cu:" WHITE_BOLD("%u") " mem:" CYAN("%zu/%zu") " MB", "OPENCL GPU",
                        device.index(),
-                       device.hasTopology() ? device.topology().toString().data() : "n/a",
+                       device.topology().toString().data(),
                        device.printableName().data(),
                        device.clock(),
                        device.computeUnits(),
@@ -144,7 +165,7 @@ public:
     }
 
 
-    inline void start()
+    inline void start(const Job &job)
     {
         LOG_INFO("%s use profile " BLUE_BG(WHITE_BOLD_S " %s ") WHITE_BOLD_S " (" CYAN_BOLD("%zu") WHITE_BOLD(" threads)") " scratchpad " CYAN_BOLD("%zu KB"),
                  tag,
@@ -161,7 +182,7 @@ public:
                        CYAN_BOLD("%3u") " |" CYAN_BOLD("%3s") " |" CYAN_BOLD("%3u") " |" CYAN("%5zu") " | %s",
                        i,
                        data.thread.index(),
-                       data.device.hasTopology() ? data.device.topology().toString().data() : "n/a",
+                       data.device.topology().toString().data(),
                        data.thread.intensity(),
                        data.thread.worksize(),
                        data.thread.stridedIndex(),
@@ -174,6 +195,8 @@ public:
                     i++;
         }
 
+        OclSharedState::start(threads, job);
+
         status.start(threads.size());
         workers.start(threads);
     }
@@ -269,7 +292,7 @@ void xmrig::OclBackend::printHashrate(bool details)
                     Hashrate::format(hashrate()->calc(i, Hashrate::MediumInterval), num + 8,     sizeof num / 3),
                     Hashrate::format(hashrate()->calc(i, Hashrate::LargeInterval),  num + 8 * 2, sizeof num / 3),
                     data.device.index(),
-                    data.device.hasTopology() ? data.device.topology().toString().data() : "n/a",
+                    data.device.topology().toString().data(),
                     data.device.printableName().data()
                     );
 
@@ -286,12 +309,15 @@ void xmrig::OclBackend::printHashrate(bool details)
 
 void xmrig::OclBackend::setJob(const Job &job)
 {
+    const OclConfig &cl = d_ptr->controller->config()->cl();
+    if (cl.isEnabled()) {
+        d_ptr->init(cl);
+    }
+
     if (!isEnabled()) {
         return stop();
     }
 
-    const OclConfig &cl = d_ptr->controller->config()->cl();
-
     std::vector<OclLaunchData> threads = cl.get(d_ptr->controller->miner(), job.algorithm(), d_ptr->platform, d_ptr->devices, tag);
     if (!d_ptr->threads.empty() && d_ptr->threads.size() == threads.size() && std::equal(d_ptr->threads.begin(), d_ptr->threads.end(), threads.begin())) {
         return;
@@ -306,7 +332,7 @@ void xmrig::OclBackend::setJob(const Job &job)
         return stop();
     }
 
-    if (!d_ptr->context.init(d_ptr->devices, threads, job)) {
+    if (!d_ptr->context.init(d_ptr->devices, threads)) {
         LOG_WARN("%s " RED_BOLD("disabled") YELLOW(" (OpenCL context unavailable)"), tag);
 
         return stop();
@@ -315,15 +341,15 @@ void xmrig::OclBackend::setJob(const Job &job)
     stop();
 
     d_ptr->threads = std::move(threads);
-    d_ptr->start();
+    d_ptr->start(job);
 }
 
 
-void xmrig::OclBackend::start(IWorker *worker)
+void xmrig::OclBackend::start(IWorker *worker, bool ready)
 {
     mutex.lock();
 
-    if (d_ptr->status.started()) {
+    if (d_ptr->status.started(ready)) {
         d_ptr->status.print();
 
         OclWorker::ready = true;
@@ -331,7 +357,9 @@ void xmrig::OclBackend::start(IWorker *worker)
 
     mutex.unlock();
 
-    worker->start();
+    if (ready) {
+        worker->start();
+    }
 }
 
 
@@ -346,6 +374,8 @@ void xmrig::OclBackend::stop()
     d_ptr->workers.stop();
     d_ptr->threads.clear();
 
+    OclSharedState::release();
+
     LOG_INFO("%s" YELLOW(" stopped") BLACK_BOLD(" (%" PRIu64 " ms)"), tag, Chrono::steadyMSecs() - ts);
 }
 
@@ -373,22 +403,19 @@ rapidjson::Value xmrig::OclBackend::toJSON(rapidjson::Document &doc) const
         return out;
     }
 
+    out.AddMember("hashrate", hashrate()->toJSON(doc), allocator);
+
     Value threads(kArrayType);
-    const Hashrate *hr = hashrate();
 
     size_t i = 0;
     for (const OclLaunchData &data : d_ptr->threads) {
         Value thread = data.thread.toJSON(doc);
         thread.AddMember("affinity", data.affinity, allocator);
+        thread.AddMember("hashrate", hashrate()->toJSON(i, doc), allocator);
 
-        Value hashrate(kArrayType);
-        hashrate.PushBack(Hashrate::normalize(hr->calc(i, Hashrate::ShortInterval)),  allocator);
-        hashrate.PushBack(Hashrate::normalize(hr->calc(i, Hashrate::MediumInterval)), allocator);
-        hashrate.PushBack(Hashrate::normalize(hr->calc(i, Hashrate::LargeInterval)),  allocator);
+        data.device.toJSON(thread, doc);
 
         i++;
-
-        thread.AddMember("hashrate", hashrate, allocator);
         threads.PushBack(thread, allocator);
     }
 

src/backend/opencl/OclBackend.h

@@ -62,7 +62,7 @@ protected:
     void prepare(const Job &nextJob) override;
     void printHashrate(bool details) override;
     void setJob(const Job &job) override;
-    void start(IWorker *worker) override;
+    void start(IWorker *worker, bool ready) override;
     void stop() override;
     void tick(uint64_t ticks) override;
 

src/backend/opencl/OclConfig.cpp

@@ -30,12 +30,16 @@
 #include "rapidjson/document.h"
 
 
+#include <algorithm>
+
+
 namespace xmrig {
 
 static const char *kAMD         = "AMD";
 static const char *kCache       = "cache";
 static const char *kCn          = "cn";
 static const char *kCn2         = "cn/2";
+static const char *kDevicesHint = "devices-hint";
 static const char *kEnabled     = "enabled";
 static const char *kINTEL       = "INTEL";
 static const char *kLoader      = "loader";
@@ -90,6 +94,22 @@ static size_t generate(const char *key, Threads<OclThreads> &threads, const Algo
 }
 
 
+static inline std::vector<OclDevice> filterDevices(const std::vector<OclDevice> &devices, const std::vector<uint32_t> &hints)
+{
+    std::vector<OclDevice> out;
+    out.reserve(std::min(devices.size(), hints.size()));
+
+    for (const auto &device  : devices) {
+        auto it = std::find(hints.begin(), hints.end(), device.index());
+        if (it != hints.end()) {
+            out.emplace_back(device);
+        }
+    }
+
+    return out;
+}
+
+
 }
 
 
@@ -173,32 +193,13 @@ std::vector<xmrig::OclLaunchData> xmrig::OclConfig::get(const Miner *miner, cons
             continue;
         }
 
-#       ifdef XMRIG_ALGO_RANDOMX
-        auto dataset = algorithm.family() == Algorithm::RANDOM_X ? std::make_shared<OclRxDataset>() : nullptr;
-#       endif
-
         if (thread.threads().size() > 1) {
-            auto interleave = std::make_shared<OclInterleave>(thread.threads().size());
-
             for (int64_t affinity : thread.threads()) {
-                OclLaunchData data(miner, algorithm, *this, platform, thread, devices[thread.index()], affinity);
-                data.interleave = interleave;
-
-#               ifdef XMRIG_ALGO_RANDOMX
-                data.dataset = dataset;
-#               endif
-
-                out.emplace_back(std::move(data));
+                out.emplace_back(miner, algorithm, *this, platform, thread, devices[thread.index()], affinity);
             }
         }
         else {
-            OclLaunchData data(miner, algorithm, *this, platform, thread, devices[thread.index()], thread.threads().front());
-
-#           ifdef XMRIG_ALGO_RANDOMX
-            data.dataset = dataset;
-#           endif
-
-            out.emplace_back(std::move(data));
+            out.emplace_back(miner, algorithm, *this, platform, thread, devices[thread.index()], thread.threads().front());
         }
     }
 
@@ -214,17 +215,20 @@ void xmrig::OclConfig::read(const rapidjson::Value &value)
         m_loader    = Json::getString(value, kLoader);
 
         setPlatform(Json::getValue(value, kPlatform));
+        setDevicesHint(Json::getString(value, kDevicesHint));
 
-        if (isEnabled()) {
-            m_threads.read(value);
+        m_threads.read(value);
 
-            generate();
-        }
+        generate();
     }
-    else if (value.IsBool() && value.IsFalse()) {
-        m_enabled = false;
+    else if (value.IsBool()) {
+        m_enabled = value.GetBool();
+
+        generate();
     }
     else {
+        m_shouldSave = true;
+
         generate();
     }
 }
@@ -232,11 +236,15 @@ void xmrig::OclConfig::read(const rapidjson::Value &value)
 
 void xmrig::OclConfig::generate()
 {
+    if (!isEnabled() || m_threads.has("*")) {
+        return;
+    }
+
     if (!OclLib::init(loader())) {
         return;
     }
 
-    const auto devices = platform().devices();
+    const auto devices = m_devicesHint.empty() ? platform().devices() : filterDevices(platform().devices(), m_devicesHint);
     if (devices.empty()) {
         return;
     }
@@ -281,6 +289,21 @@ void xmrig::OclConfig::generate()
 }
 
 
+void xmrig::OclConfig::setDevicesHint(const char *devicesHint)
+{
+    if (devicesHint == nullptr) {
+        return;
+    }
+
+    const auto indexes = String(devicesHint).split(',');
+    m_devicesHint.reserve(indexes.size());
+
+    for (const auto &index : indexes) {
+        m_devicesHint.push_back(strtoul(index, nullptr, 10));
+    }
+}
+
+
 void xmrig::OclConfig::setPlatform(const rapidjson::Value &platform)
 {
     if (platform.IsString()) {

src/backend/opencl/OclConfig.h

@@ -53,11 +53,13 @@ public:
 
 private:
     void generate();
+    void setDevicesHint(const char *devicesHint);
     void setPlatform(const rapidjson::Value &platform);
 
     bool m_cache         = true;
-    bool m_enabled       = true;
+    bool m_enabled       = false;
     bool m_shouldSave    = false;
+    std::vector<uint32_t> m_devicesHint;
     String m_loader;
     String m_platformVendor;
     Threads<OclThreads> m_threads;

src/backend/opencl/OclLaunchData.cpp

@@ -25,6 +25,8 @@
 
 
 #include "backend/opencl/OclLaunchData.h"
+
+#include "backend/common/Tags.h"
 #include "backend/opencl/OclConfig.h"
 
 
@@ -45,3 +47,9 @@ bool xmrig::OclLaunchData::isEqual(const OclLaunchData &other) const
     return (other.algorithm == algorithm &&
             other.thread    == thread);
 }
+
+
+const char *xmrig::OclLaunchData::tag()
+{
+    return ocl_tag();
+}

src/backend/opencl/OclLaunchData.h

@@ -27,19 +27,14 @@
 #define XMRIG_OCLLAUNCHDATA_H
 
 
-#include "backend/opencl/OclInterleave.h"
 #include "backend/opencl/OclThread.h"
+#include "backend/opencl/runners/tools/OclSharedData.h"
 #include "backend/opencl/wrappers/OclDevice.h"
 #include "backend/opencl/wrappers/OclPlatform.h"
 #include "crypto/common/Algorithm.h"
 #include "crypto/common/Nonce.h"
 
 
-#ifdef XMRIG_ALGO_RANDOMX
-#   include "backend/opencl/runners/tools/OclRxDataset.h"
-#endif
-
-
 using cl_context = struct _cl_context *;
 
 
@@ -62,6 +57,8 @@ public:
     inline bool operator!=(const OclLaunchData &other) const    { return !isEqual(other); }
     inline bool operator==(const OclLaunchData &other) const    { return isEqual(other); }
 
+    static const char *tag();
+
     cl_context ctx = nullptr;
     const Algorithm algorithm;
     const bool cache;
@@ -70,11 +67,6 @@ public:
     const OclDevice device;
     const OclPlatform platform;
     const OclThread thread;
-    OclInterleavePtr interleave;
-
-#   ifdef XMRIG_ALGO_RANDOMX
-    OclRxDatasetPtr dataset;
-#   endif
 };
 
 

src/backend/opencl/OclThread.cpp

@@ -52,6 +52,10 @@ static const char* kDatasetHost  = "dataset_host";
 
 xmrig::OclThread::OclThread(const rapidjson::Value &value)
 {
+    if (!value.IsObject()) {
+        return;
+    }
+
     m_index         = Json::getUint(value, kIndex);
     m_worksize      = std::max(std::min(Json::getUint(value, kWorksize), 128u), 1u);
     m_unrollFactor  = std::max(std::min(Json::getUint(value, kUnroll, m_unrollFactor), 128u), 1u);

src/backend/opencl/OclWorker.cpp

@@ -28,6 +28,8 @@
 
 #include "backend/common/Tags.h"
 #include "backend/opencl/runners/OclCnRunner.h"
+#include "backend/opencl/runners/tools/OclSharedData.h"
+#include "backend/opencl/runners/tools/OclSharedState.h"
 #include "base/io/log/Log.h"
 #include "base/tools/Chrono.h"
 #include "core/Miner.h"
@@ -75,7 +77,7 @@ xmrig::OclWorker::OclWorker(size_t id, const OclLaunchData &data) :
     m_algorithm(data.algorithm),
     m_miner(data.miner),
     m_intensity(data.thread.intensity()),
-    m_interleave(data.interleave)
+    m_sharedData(OclSharedState::get(data.device.index()))
 {
     switch (m_algorithm.family()) {
     case Algorithm::RANDOM_X:
@@ -137,15 +139,19 @@ bool xmrig::OclWorker::selfTest()
 }
 
 
+size_t xmrig::OclWorker::intensity() const
+{
+    return m_runner ? m_runner->intensity() : 0;
+}
+
+
 void xmrig::OclWorker::start()
 {
     cl_uint results[0x100];
 
     while (Nonce::sequence(Nonce::OPENCL) > 0) {
         if (!isReady()) {
-            if (m_interleave) {
-                m_interleave->setResumeCounter(0);
-            }
+            m_sharedData.setResumeCounter(0);
 
             do {
                 std::this_thread::sleep_for(std::chrono::milliseconds(200));
@@ -156,9 +162,7 @@ void xmrig::OclWorker::start()
                 break;
             }
 
-            if (m_interleave) {
-                m_interleave->resumeDelay(m_id);
-            }
+            m_sharedData.resumeDelay(m_id);
 
             if (!consumeJob()) {
                 return;
@@ -166,9 +170,7 @@ void xmrig::OclWorker::start()
         }
 
         while (!Nonce::isOutdated(Nonce::OPENCL, m_job.sequence())) {
-            if (m_interleave) {
-                m_interleave->adjustDelay(m_id);
-            }
+            m_sharedData.adjustDelay(m_id);
 
             const uint64_t t = Chrono::steadyMSecs();
 
@@ -227,9 +229,7 @@ void xmrig::OclWorker::storeStats(uint64_t t)
 
     m_count += m_intensity;
 
-    if (m_interleave) {
-        m_interleave->setRunTime(Chrono::steadyMSecs() - t);
-    }
+    m_sharedData.setRunTime(Chrono::steadyMSecs() - t);
 
     Worker::storeStats();
 }

src/backend/opencl/OclWorker.h

@@ -56,6 +56,7 @@ public:
 
 protected:
     bool selfTest() override;
+    size_t intensity() const override;
     void start() override;
 
 private:
@@ -66,7 +67,7 @@ private:
     const Miner *m_miner;
     const uint32_t m_intensity;
     IOclRunner *m_runner = nullptr;
-    OclInterleavePtr m_interleave;
+    OclSharedData &m_sharedData;
     WorkerJob<1> m_job;
 };
 

src/backend/opencl/cl/rx/aes.cl

@@ -604,7 +604,7 @@ __kernel void hashAes1Rx4(__global const void* input, __global void* hash, uint
 	__local const uint* const t2 = ((sub & 1) == 0) ? (T + 512) : (T + 1536);
 	__local const uint* const t3 = ((sub & 1) == 0) ? (T + 768) : (T + 1280);
 
-	#pragma unroll(8)
+	#pragma unroll 8
 	for (uint i = 0; i < inputSize / sizeof(uint4); i += 4, p += 4)
 	{
 		uint k[4], y[4];

src/backend/opencl/cl/rx/fillAes1Rx4.cl

@@ -72,7 +72,7 @@ __kernel void fillAes_name(__global void* state, __global void* out, uint batch_
 	const __local uint* const t2 = (sub & 1) ? (T + 512) : (T + 1536);
 	const __local uint* const t3 = (sub & 1) ? (T + 768) : (T + 1280);
 
-	#pragma unroll(unroll_factor)
+	#pragma unroll unroll_factor
 	for (uint i = 0; i < outputSize / sizeof(uint4); i += 4, p += 4)
 	{
 		uint y[4];

src/backend/opencl/cl/rx/randomx_jit.cl

@@ -902,7 +902,7 @@ __global uint* generate_jit_code(__global uint2* e, __global uint2* p0, __global
 {
 	int prefetch_data_count;
 
-	#pragma unroll(1)
+	#pragma unroll 1
 	for (int pass = 0; pass < 2; ++pass)
 	{
 #if RANDOMX_PROGRAM_SIZE > 256
@@ -929,7 +929,7 @@ __global uint* generate_jit_code(__global uint2* e, __global uint2* p0, __global
 
 		prefetch_data_count = 0;
 
-		#pragma unroll(1)
+		#pragma unroll 1
 		for (uint i = 0; i < RANDOMX_PROGRAM_SIZE; ++i)
 		{
 			// Clean flags
@@ -1314,7 +1314,7 @@ __global uint* generate_jit_code(__global uint2* e, __global uint2* p0, __global
 
 	// Sort p0
 	uint prev = p0[0].x;
-	#pragma unroll(1)
+	#pragma unroll 1
 	for (int j = 1; j < prefetch_data_count; ++j)
 	{
 		uint2 cur = p0[j];
@@ -1344,7 +1344,7 @@ __global uint* generate_jit_code(__global uint2* e, __global uint2* p0, __global
 
 	__global int* prefetched_vgprs = prefecth_vgprs_stack + num_prefetch_vgprs;
 
-	#pragma unroll(8)
+	#pragma unroll 8
 	for (int i = 0; i < RANDOMX_PROGRAM_SIZE; ++i)
 		prefetched_vgprs[i] = 0;
 
@@ -1359,7 +1359,7 @@ __global uint* generate_jit_code(__global uint2* e, __global uint2* p0, __global
 	const uint size_limit = (COMPILED_PROGRAM_SIZE - 200) / sizeof(uint);
 	__global uint* start_p = p;
 
-	#pragma unroll(1)
+	#pragma unroll 1
 	for (int i = 0; i < RANDOMX_PROGRAM_SIZE; ++i)
 	{
 		const uint2 inst = e[i];

src/backend/opencl/cl/rx/randomx_vm.cl

@@ -1658,7 +1658,7 @@ uint32_t inner_loop(
 	const int32_t sub2 = sub >> 1;
 	imm_buf[IMM_INDEX_COUNT + 1] = fprc;
 
-	#pragma unroll(1)
+	#pragma unroll 1
 	for (int32_t ip = 0; ip < program_length;)
 	{
 		imm_buf[IMM_INDEX_COUNT] = ip;
@@ -1934,7 +1934,7 @@ __kernel void execute_vm(__global void* vm_states, __global void* rounding, __gl
 	const uint32_t workers_mask = ((1 << WORKERS_PER_HASH) - 1) << ((get_local_id(0) / IDX_WIDTH) * IDX_WIDTH);
 	const uint32_t fp_workers_mask = 3 << (((sub >> 1) << 1) + (get_local_id(0) / IDX_WIDTH) * IDX_WIDTH);
 
-	#pragma unroll(1)
+	#pragma unroll 1
 	for (int ic = 0; ic < num_iterations; ++ic)
 	{
 		__local uint64_t *r;

src/backend/opencl/generators/ocl_generic_cn_generator.cpp

@@ -45,7 +45,13 @@ static inline uint32_t getMaxThreads(const OclDevice &device, const Algorithm &a
         return 40000u;
     }
 
-    return ((algorithm.l3() <= oneMiB) ? 2u : 1u) * 1000u;
+    const uint32_t ratio = (algorithm.l3() <= oneMiB) ? 2u : 1u;
+
+    if (device.vendorId() == OCL_VENDOR_INTEL) {
+        return ratio * device.computeUnits() * 8;
+    }
+
+    return ratio * 1000u;
 }
 
 
@@ -107,7 +113,7 @@ bool ocl_generic_cn_generator(const OclDevice &device, const Algorithm &algorith
         return false;
     }
 
-    const uint32_t threadCount = ((device.globalMemSize() - intensity * 2 * algorithm.l3()) > 128 * oneMiB) ? 2 : 1;
+    const uint32_t threadCount = (device.vendorId() == OCL_VENDOR_AMD && (device.globalMemSize() - intensity * 2 * algorithm.l3()) > 128 * oneMiB) ? 2 : 1;
 
     threads.add(OclThread(device.index(), intensity, 8, getStridedIndex(device, algorithm), 2, threadCount, 8));
 

src/backend/opencl/generators/ocl_generic_rx_generator.cpp

@@ -86,7 +86,7 @@ bool ocl_generic_rx_generator(const OclDevice &device, const Algorithm &algorith
         return false;
     }
 
-    threads.add(OclThread(device.index(), intensity, 8, 2, gcnAsm, datasetHost, 6));
+    threads.add(OclThread(device.index(), intensity, 8, device.vendorId() == OCL_VENDOR_AMD ? 2 : 1, gcnAsm, datasetHost, 6));
 
     return true;
 }

src/backend/opencl/interfaces/IOclRunner.h

@@ -51,18 +51,19 @@ public:
     IOclRunner()          = default;
     virtual ~IOclRunner() = default;
 
-    virtual void run(uint32_t nonce, uint32_t *hashOutput)  = 0;
-    virtual void set(const Job &job, uint8_t *blob)         = 0;
     virtual cl_context ctx() const                          = 0;
     virtual const Algorithm &algorithm() const              = 0;
     virtual const char *buildOptions() const                = 0;
     virtual const char *deviceKey() const                   = 0;
     virtual const char *source() const                      = 0;
     virtual const OclLaunchData &data() const               = 0;
+    virtual size_t intensity() const                        = 0;
     virtual size_t threadId() const                         = 0;
     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 set(const Job &job, uint8_t *blob)         = 0;
 
 protected:
     virtual size_t bufferSize() const                       = 0;

src/backend/opencl/opencl.cmake

@@ -14,7 +14,6 @@ if (WITH_OPENCL)
         src/backend/opencl/OclCache.h
         src/backend/opencl/OclConfig.h
         src/backend/opencl/OclGenerator.h
-        src/backend/opencl/OclInterleave.h
         src/backend/opencl/OclLaunchData.h
         src/backend/opencl/OclThread.h
         src/backend/opencl/OclThreads.h
@@ -22,6 +21,8 @@ if (WITH_OPENCL)
         src/backend/opencl/runners/OclBaseRunner.h
         src/backend/opencl/runners/OclCnRunner.h
         src/backend/opencl/runners/tools/OclCnR.h
+        src/backend/opencl/runners/tools/OclSharedData.h
+        src/backend/opencl/runners/tools/OclSharedState.h
         src/backend/opencl/wrappers/OclContext.h
         src/backend/opencl/wrappers/OclDevice.h
         src/backend/opencl/wrappers/OclError.h
@@ -42,7 +43,6 @@ if (WITH_OPENCL)
         src/backend/opencl/OclBackend.cpp
         src/backend/opencl/OclCache.cpp
         src/backend/opencl/OclConfig.cpp
-        src/backend/opencl/OclInterleave.cpp
         src/backend/opencl/OclLaunchData.cpp
         src/backend/opencl/OclThread.cpp
         src/backend/opencl/OclThreads.cpp
@@ -50,6 +50,8 @@ if (WITH_OPENCL)
         src/backend/opencl/runners/OclBaseRunner.cpp
         src/backend/opencl/runners/OclCnRunner.cpp
         src/backend/opencl/runners/tools/OclCnR.cpp
+        src/backend/opencl/runners/tools/OclSharedData.cpp
+        src/backend/opencl/runners/tools/OclSharedState.cpp
         src/backend/opencl/wrappers/OclContext.cpp
         src/backend/opencl/wrappers/OclDevice.cpp
         src/backend/opencl/wrappers/OclError.cpp
@@ -78,7 +80,6 @@ if (WITH_OPENCL)
            src/backend/opencl/runners/OclRxBaseRunner.h
            src/backend/opencl/runners/OclRxJitRunner.h
            src/backend/opencl/runners/OclRxVmRunner.h
-           src/backend/opencl/runners/tools/OclRxDataset.h
            )
 
        list(APPEND SOURCES_BACKEND_OPENCL
@@ -95,7 +96,6 @@ if (WITH_OPENCL)
            src/backend/opencl/runners/OclRxBaseRunner.cpp
            src/backend/opencl/runners/OclRxJitRunner.cpp
            src/backend/opencl/runners/OclRxVmRunner.cpp
-           src/backend/opencl/runners/tools/OclRxDataset.cpp
            )
    endif()
 

src/backend/opencl/runners/OclBaseRunner.cpp

@@ -23,10 +23,11 @@
  */
 
 
+#include "backend/opencl/runners/OclBaseRunner.h"
 #include "backend/opencl/cl/OclSource.h"
 #include "backend/opencl/OclCache.h"
 #include "backend/opencl/OclLaunchData.h"
-#include "backend/opencl/runners/OclBaseRunner.h"
+#include "backend/opencl/runners/tools/OclSharedState.h"
 #include "backend/opencl/wrappers/OclError.h"
 #include "backend/opencl/wrappers/OclLib.h"
 #include "base/io/log/Log.h"
@@ -34,13 +35,17 @@
 #include "crypto/common/VirtualMemory.h"
 
 
+constexpr size_t oneGiB = 1024 * 1024 * 1024;
+
+
 xmrig::OclBaseRunner::OclBaseRunner(size_t id, const OclLaunchData &data) :
     m_algorithm(data.algorithm),
     m_ctx(data.ctx),
     m_source(OclSource::get(data.algorithm)),
     m_data(data),
     m_align(OclLib::getUint(data.device.id(), CL_DEVICE_MEM_BASE_ADDR_ALIGN)),
-    m_threadId(id)
+    m_threadId(id),
+    m_intensity(data.thread.intensity())
 {
     m_deviceKey = data.device.name();
 
@@ -92,16 +97,17 @@ void xmrig::OclBaseRunner::build()
 
 void xmrig::OclBaseRunner::init()
 {
-    m_queue  = OclLib::createCommandQueue(m_ctx, data().device.id());
+    m_queue = OclLib::createCommandQueue(m_ctx, data().device.id());
 
-    constexpr size_t oneGiB = 1024 * 1024 * 1024;
-    size_t size             = bufferSize();
+    size_t size = align(bufferSize());
 
-    if (size < oneGiB && data().device.freeMemSize() >= oneGiB) {
-        size = oneGiB;
+    if (size < oneGiB && data().device.vendorId() == OCL_VENDOR_AMD && data().device.freeMemSize() >= oneGiB) {
+        m_buffer = OclSharedState::get(data().device.index()).createBuffer(m_ctx, size, m_offset);
+    }
+    else {
+        m_buffer = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, size);
     }
 
-    m_buffer = OclLib::createBuffer(m_ctx, CL_MEM_READ_WRITE, size);
     m_input  = createSubBuffer(CL_MEM_READ_ONLY | CL_MEM_HOST_WRITE_ONLY, Job::kMaxBlobSize);
     m_output = createSubBuffer(CL_MEM_READ_WRITE, sizeof(cl_uint) * 0x100);
 }

src/backend/opencl/runners/OclBaseRunner.h

@@ -55,6 +55,7 @@ protected:
     inline const char *deviceKey() const override       { return m_deviceKey.c_str(); }
     inline const char *source() const override          { return m_source; }
     inline const OclLaunchData &data() const override   { return m_data; }
+    inline size_t intensity() const override            { return m_intensity; }
     inline size_t threadId() const override             { return m_threadId; }
 
     size_t bufferSize() const override;
@@ -83,6 +84,7 @@ protected:
     size_t m_offset             = 0;
     std::string m_deviceKey;
     std::string m_options;
+    uint32_t m_intensity;
 };
 
 

src/backend/opencl/runners/OclCnRunner.cpp

@@ -83,12 +83,10 @@ xmrig::OclCnRunner::~OclCnRunner()
 
 size_t xmrig::OclCnRunner::bufferSize() const
 {
-    const size_t g_thd = data().thread.intensity();
-
     return OclBaseRunner::bufferSize() +
-           align(m_algorithm.l3() * g_thd) +
-           align(200 * g_thd) +
-           (align(sizeof(cl_uint) * (g_thd + 2)) * BRANCH_MAX);
+           align(m_algorithm.l3() * m_intensity) +
+           align(200 * m_intensity) +
+           (align(sizeof(cl_uint) * (m_intensity + 2)) * BRANCH_MAX);
 }
 
 
@@ -96,14 +94,13 @@ void xmrig::OclCnRunner::run(uint32_t nonce, uint32_t *hashOutput)
 {
     static const cl_uint zero = 0;
 
-    const size_t g_intensity = data().thread.intensity();
-    const size_t w_size      = data().thread.worksize();
-    const size_t g_thd       = ((g_intensity + w_size - 1u) / w_size) * w_size;
+    const size_t w_size = data().thread.worksize();
+    const size_t g_thd  = ((m_intensity + w_size - 1u) / w_size) * w_size;
 
     assert(g_thd % w_size == 0);
 
     for (size_t i = 0; i < BRANCH_MAX; ++i) {
-        enqueueWriteBuffer(m_branches[i], CL_FALSE, sizeof(cl_uint) * g_intensity, sizeof(cl_uint), &zero);
+        enqueueWriteBuffer(m_branches[i], CL_FALSE, sizeof(cl_uint) * m_intensity, sizeof(cl_uint), &zero);
     }
 
     enqueueWriteBuffer(m_output, CL_FALSE, sizeof(cl_uint) * 0xFF, sizeof(cl_uint), &zero);
@@ -134,10 +131,15 @@ void xmrig::OclCnRunner::set(const Job &job, uint8_t *blob)
     if (m_algorithm == Algorithm::CN_R && m_height != job.height()) {
         delete m_cn1;
 
-        m_height = job.height();
-        m_cnr    = OclCnR::get(*this, m_height);
-        m_cn1    = new Cn1Kernel(m_cnr, m_height);
-        m_cn1->setArgs(m_input, m_scratchpads, m_states, data().thread.intensity());
+        m_height     = job.height();
+        auto program = OclCnR::get(*this, m_height);
+        m_cn1        = new Cn1Kernel(program, m_height);
+        m_cn1->setArgs(m_input, m_scratchpads, m_states, m_intensity);
+
+        if (m_cnr != program) {
+            OclLib::release(m_cnr);
+            m_cnr = OclLib::retain(program);
+        }
     }
 
     for (auto kernel : m_branchKernels) {
@@ -150,22 +152,20 @@ void xmrig::OclCnRunner::build()
 {
     OclBaseRunner::build();
 
-    const uint32_t intensity = data().thread.intensity();
-
     m_cn0 = new Cn0Kernel(m_program);
-    m_cn0->setArgs(m_input, m_scratchpads, m_states, intensity);
+    m_cn0->setArgs(m_input, m_scratchpads, m_states, m_intensity);
 
     m_cn2 = new Cn2Kernel(m_program);
-    m_cn2->setArgs(m_scratchpads, m_states, m_branches, intensity);
+    m_cn2->setArgs(m_scratchpads, m_states, m_branches, m_intensity);
 
     if (m_algorithm != Algorithm::CN_R) {
         m_cn1 = new Cn1Kernel(m_program);
-        m_cn1->setArgs(m_input, m_scratchpads, m_states, intensity);
+        m_cn1->setArgs(m_input, m_scratchpads, m_states, m_intensity);
     }
 
     for (size_t i = 0; i < BRANCH_MAX; ++i) {
         auto kernel = new CnBranchKernel(i, m_program);
-        kernel->setArgs(m_states, m_branches[i], m_output, intensity);
+        kernel->setArgs(m_states, m_branches[i], m_output, m_intensity);
 
         m_branchKernels[i] = kernel;
     }
@@ -176,12 +176,10 @@ void xmrig::OclCnRunner::init()
 {
     OclBaseRunner::init();
 
-    const size_t g_thd = data().thread.intensity();
-
-    m_scratchpads = createSubBuffer(CL_MEM_READ_WRITE, m_algorithm.l3() * g_thd);
-    m_states      = createSubBuffer(CL_MEM_READ_WRITE, 200 * g_thd);
+    m_scratchpads = createSubBuffer(CL_MEM_READ_WRITE, m_algorithm.l3() * m_intensity);
+    m_states      = createSubBuffer(CL_MEM_READ_WRITE, 200 * m_intensity);
 
     for (size_t i = 0; i < BRANCH_MAX; ++i) {
-        m_branches[i] = createSubBuffer(CL_MEM_READ_WRITE, sizeof(cl_uint) * (g_thd + 2));
+        m_branches[i] = createSubBuffer(CL_MEM_READ_WRITE, sizeof(cl_uint) * (m_intensity + 2));
     }
 }

src/backend/opencl/runners/OclRxBaseRunner.cpp

@@ -30,6 +30,7 @@
 #include "backend/opencl/kernels/rx/FindSharesKernel.h"
 #include "backend/opencl/kernels/rx/HashAesKernel.h"
 #include "backend/opencl/OclLaunchData.h"
+#include "backend/opencl/runners/tools/OclSharedState.h"
 #include "backend/opencl/wrappers/OclLib.h"
 #include "base/net/stratum/Job.h"
 #include "crypto/rx/Rx.h"
@@ -75,6 +76,7 @@ xmrig::OclRxBaseRunner::~OclRxBaseRunner()
     OclLib::release(m_hashes);
     OclLib::release(m_rounding);
     OclLib::release(m_scratchpads);
+    OclLib::release(m_dataset);
 }
 
 
@@ -87,28 +89,26 @@ void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
 
     enqueueWriteBuffer(m_output, CL_FALSE, sizeof(cl_uint) * 0xFF, sizeof(uint32_t), &zero);
 
-    const uint32_t g_intensity = data().thread.intensity();
-
-    m_blake2b_initial_hash->enqueue(m_queue, g_intensity);
-    m_fillAes1Rx4_scratchpad->enqueue(m_queue, g_intensity);
+    m_blake2b_initial_hash->enqueue(m_queue, m_intensity);
+    m_fillAes1Rx4_scratchpad->enqueue(m_queue, m_intensity);
 
     const uint32_t programCount = RxAlgo::programCount(m_algorithm);
 
     for (uint32_t i = 0; i < programCount; ++i) {
-        m_fillAes4Rx4_entropy->enqueue(m_queue, g_intensity);
+        m_fillAes4Rx4_entropy->enqueue(m_queue, m_intensity);
 
         execute(i);
 
         if (i == programCount - 1) {
-            m_hashAes1Rx4->enqueue(m_queue, g_intensity);
-            m_blake2b_hash_registers_32->enqueue(m_queue, g_intensity);
+            m_hashAes1Rx4->enqueue(m_queue, m_intensity);
+            m_blake2b_hash_registers_32->enqueue(m_queue, m_intensity);
         }
         else {
-            m_blake2b_hash_registers_64->enqueue(m_queue, g_intensity);
+            m_blake2b_hash_registers_64->enqueue(m_queue, m_intensity);
         }
     }
 
-    m_find_shares->enqueue(m_queue, g_intensity);
+    m_find_shares->enqueue(m_queue, m_intensity);
 
     finalize(hashOutput);
 
@@ -122,7 +122,7 @@ void xmrig::OclRxBaseRunner::set(const Job &job, uint8_t *blob)
         m_seed = job.seed();
 
         auto dataset = Rx::dataset(job, 0);
-        enqueueWriteBuffer(data().dataset->get(), CL_TRUE, 0, dataset->size(), dataset->raw());
+        enqueueWriteBuffer(m_dataset, CL_TRUE, 0, dataset->size(), dataset->raw());
     }
 
     if (job.size() < Job::kMaxBlobSize) {
@@ -138,13 +138,11 @@ void xmrig::OclRxBaseRunner::set(const Job &job, uint8_t *blob)
 
 size_t xmrig::OclRxBaseRunner::bufferSize() const
 {
-    const size_t g_thd = data().thread.intensity();
-
     return OclBaseRunner::bufferSize() +
-           align((m_algorithm.l3() + 64) * g_thd) +
-           align(64 * g_thd) +
-           align((128 + 2560) * g_thd) +
-           align(sizeof(uint32_t) * g_thd);
+           align((m_algorithm.l3() + 64) * m_intensity) +
+           align(64 * m_intensity) +
+           align((128 + 2560) * m_intensity) +
+           align(sizeof(uint32_t) * m_intensity);
 }
 
 
@@ -152,14 +150,13 @@ void xmrig::OclRxBaseRunner::build()
 {
     OclBaseRunner::build();
 
-    const uint32_t batch_size = data().thread.intensity();
     const uint32_t rx_version = RxAlgo::version(m_algorithm);
 
     m_fillAes1Rx4_scratchpad = new FillAesKernel(m_program, "fillAes1Rx4_scratchpad");
-    m_fillAes1Rx4_scratchpad->setArgs(m_hashes, m_scratchpads, batch_size, rx_version);
+    m_fillAes1Rx4_scratchpad->setArgs(m_hashes, m_scratchpads, m_intensity, rx_version);
 
     m_fillAes4Rx4_entropy = new FillAesKernel(m_program, "fillAes4Rx4_entropy");
-    m_fillAes4Rx4_entropy->setArgs(m_hashes, m_entropy, batch_size, rx_version);
+    m_fillAes4Rx4_entropy->setArgs(m_hashes, m_entropy, m_intensity, rx_version);
 
     m_hashAes1Rx4 = new HashAesKernel(m_program);
 
@@ -178,10 +175,9 @@ void xmrig::OclRxBaseRunner::init()
 {
     OclBaseRunner::init();
 
-    const size_t g_thd = data().thread.intensity();
-
-    m_scratchpads = createSubBuffer(CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, (m_algorithm.l3() + 64) * g_thd);
-    m_hashes      = createSubBuffer(CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, 64 * g_thd);
-    m_entropy     = createSubBuffer(CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, (128 + 2560) * g_thd);
-    m_rounding    = createSubBuffer(CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, sizeof(uint32_t) * g_thd);
+    m_scratchpads = createSubBuffer(CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, (m_algorithm.l3() + 64) * m_intensity);
+    m_hashes      = createSubBuffer(CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, 64 * m_intensity);
+    m_entropy     = createSubBuffer(CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, (128 + 2560) * m_intensity);
+    m_rounding    = createSubBuffer(CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, sizeof(uint32_t) * m_intensity);
+    m_dataset     = OclSharedState::get(data().device.index()).dataset();
 }

src/backend/opencl/runners/OclRxBaseRunner.h

@@ -62,6 +62,7 @@ protected:
     Blake2bHashRegistersKernel *m_blake2b_hash_registers_64 = nullptr;
     Blake2bInitialHashKernel *m_blake2b_initial_hash        = nullptr;
     Buffer m_seed;
+    cl_mem m_dataset                                        = nullptr;
     cl_mem m_entropy                                        = nullptr;
     cl_mem m_hashes                                         = nullptr;
     cl_mem m_rounding                                       = nullptr;

src/backend/opencl/runners/OclRxJitRunner.cpp

@@ -54,9 +54,7 @@ xmrig::OclRxJitRunner::~OclRxJitRunner()
 
 size_t xmrig::OclRxJitRunner::bufferSize() const
 {
-    const size_t g_thd = data().thread.intensity();
-
-    return OclRxBaseRunner::bufferSize() + align(256 * g_thd) + align(5120 * g_thd) + align(10048 * g_thd);
+    return OclRxBaseRunner::bufferSize() + align(256 * m_intensity) + align(5120 * m_intensity) + align(10048 * m_intensity);
 }
 
 
@@ -64,33 +62,29 @@ void xmrig::OclRxJitRunner::build()
 {
     OclRxBaseRunner::build();
 
-    const uint32_t batch_size = data().thread.intensity();
-
-    m_hashAes1Rx4->setArgs(m_scratchpads, m_registers, 256, batch_size);
+    m_hashAes1Rx4->setArgs(m_scratchpads, m_registers, 256, m_intensity);
     m_blake2b_hash_registers_32->setArgs(m_hashes, m_registers, 256);
     m_blake2b_hash_registers_64->setArgs(m_hashes, m_registers, 256);
 
     m_randomx_jit = new RxJitKernel(m_program);
-    m_randomx_jit->setArgs(m_entropy, m_registers, m_intermediate_programs, m_programs, batch_size, m_rounding);
+    m_randomx_jit->setArgs(m_entropy, m_registers, m_intermediate_programs, m_programs, m_intensity, m_rounding);
 
     if (!loadAsmProgram()) {
         throw std::runtime_error(OclError::toString(CL_INVALID_PROGRAM));
     }
 
     m_randomx_run = new RxRunKernel(m_asmProgram);
-    m_randomx_run->setArgs(data().dataset->get(), m_scratchpads, m_registers, m_rounding, m_programs, batch_size, m_algorithm);
+    m_randomx_run->setArgs(m_dataset, m_scratchpads, m_registers, m_rounding, m_programs, m_intensity, m_algorithm);
 }
 
 
 void xmrig::OclRxJitRunner::execute(uint32_t iteration)
 {
-    const uint32_t g_intensity = data().thread.intensity();
-
-    m_randomx_jit->enqueue(m_queue, g_intensity, iteration);
+    m_randomx_jit->enqueue(m_queue, m_intensity, iteration);
 
     OclLib::finish(m_queue);
 
-    m_randomx_run->enqueue(m_queue, g_intensity);
+    m_randomx_run->enqueue(m_queue, m_intensity);
 }
 
 
@@ -98,11 +92,9 @@ void xmrig::OclRxJitRunner::init()
 {
     OclRxBaseRunner::init();
 
-    const size_t g_thd = data().thread.intensity();
-
-    m_registers             = createSubBuffer(CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, 256 * g_thd);
-    m_intermediate_programs = createSubBuffer(CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, 5120 * g_thd);
-    m_programs              = createSubBuffer(CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, 10048 * g_thd);
+    m_registers             = createSubBuffer(CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, 256 * m_intensity);
+    m_intermediate_programs = createSubBuffer(CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, 5120 * m_intensity);
+    m_programs              = createSubBuffer(CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, 10048 * m_intensity);
 }
 
 

src/backend/opencl/runners/OclRxVmRunner.cpp

@@ -52,7 +52,7 @@ xmrig::OclRxVmRunner::~OclRxVmRunner()
 
 size_t xmrig::OclRxVmRunner::bufferSize() const
 {
-    return OclRxBaseRunner::bufferSize() + (align(2560 * data().thread.intensity()));
+    return OclRxBaseRunner::bufferSize() + (align(2560 * m_intensity));
 }
 
 
@@ -60,10 +60,9 @@ void xmrig::OclRxVmRunner::build()
 {
     OclRxBaseRunner::build();
 
-    const uint32_t batch_size       = data().thread.intensity();
-    const uint32_t hashStrideBytes  = RxAlgo::programSize(m_algorithm) * 8;
+    const uint32_t hashStrideBytes = RxAlgo::programSize(m_algorithm) * 8;
 
-    m_hashAes1Rx4->setArgs(m_scratchpads, m_vm_states, hashStrideBytes, batch_size);
+    m_hashAes1Rx4->setArgs(m_scratchpads, m_vm_states, hashStrideBytes, m_intensity);
     m_blake2b_hash_registers_32->setArgs(m_hashes, m_vm_states, hashStrideBytes);
     m_blake2b_hash_registers_64->setArgs(m_hashes, m_vm_states, hashStrideBytes);
 
@@ -71,7 +70,7 @@ void xmrig::OclRxVmRunner::build()
     m_init_vm->setArgs(m_entropy, m_vm_states, m_rounding);
 
     m_execute_vm = new ExecuteVmKernel(m_program);
-    m_execute_vm->setArgs(m_vm_states, m_rounding, m_scratchpads, data().dataset->get(), batch_size);
+    m_execute_vm->setArgs(m_vm_states, m_rounding, m_scratchpads, m_dataset, m_intensity);
 }
 
 
@@ -79,9 +78,8 @@ void xmrig::OclRxVmRunner::execute(uint32_t iteration)
 {
     const uint32_t bfactor        = std::min(data().thread.bfactor(), 8u);
     const uint32_t num_iterations = RxAlgo::programIterations(m_algorithm) >> bfactor;
-    const uint32_t g_intensity    = data().thread.intensity();
 
-    m_init_vm->enqueue(m_queue, g_intensity, iteration);
+    m_init_vm->enqueue(m_queue, m_intensity, iteration);
 
     m_execute_vm->setIterations(num_iterations);
 
@@ -90,7 +88,7 @@ void xmrig::OclRxVmRunner::execute(uint32_t iteration)
             m_execute_vm->setLast(1);
         }
 
-        m_execute_vm->enqueue(m_queue, g_intensity, m_worksize);
+        m_execute_vm->enqueue(m_queue, m_intensity, m_worksize);
 
         if (j == 0) {
             m_execute_vm->setFirst(0);
@@ -103,5 +101,5 @@ void xmrig::OclRxVmRunner::init()
 {
     OclRxBaseRunner::init();
 
-    m_vm_states = createSubBuffer(CL_MEM_READ_WRITE, 2560 * data().thread.intensity());
+    m_vm_states = createSubBuffer(CL_MEM_READ_WRITE, 2560 * m_intensity);
 }

src/backend/opencl/runners/OclRyoRunner.cpp

@@ -63,9 +63,7 @@ xmrig::OclRyoRunner::~OclRyoRunner()
 
 size_t xmrig::OclRyoRunner::bufferSize() const
 {
-    const size_t g_thd = data().thread.intensity();
-
-    return OclBaseRunner::bufferSize() + align(data().algorithm.l3() * g_thd) + align(200 * g_thd);
+    return OclBaseRunner::bufferSize() + align(data().algorithm.l3() * m_intensity) + align(200 * m_intensity);
 }
 
 
@@ -73,9 +71,8 @@ void xmrig::OclRyoRunner::run(uint32_t nonce, uint32_t *hashOutput)
 {
     static const cl_uint zero = 0;
 
-    const size_t g_intensity = data().thread.intensity();
-    const size_t w_size      = data().thread.worksize();
-    const size_t g_thd       = ((g_intensity + w_size - 1u) / w_size) * w_size;
+    const size_t w_size = data().thread.worksize();
+    const size_t g_thd  = ((m_intensity + w_size - 1u) / w_size) * w_size;
 
     assert(g_thd % w_size == 0);
 
@@ -109,19 +106,17 @@ void xmrig::OclRyoRunner::build()
 {
     OclBaseRunner::build();
 
-    const uint32_t intensity = data().thread.intensity();
-
     m_cn00 = new Cn00RyoKernel(m_program);
     m_cn00->setArgs(m_scratchpads, m_states);
 
     m_cn0 = new Cn0Kernel(m_program);
-    m_cn0->setArgs(m_input, m_scratchpads, m_states, intensity);
+    m_cn0->setArgs(m_input, m_scratchpads, m_states, m_intensity);
 
     m_cn1 = new Cn1RyoKernel(m_program);
-    m_cn1->setArgs(m_scratchpads, m_states, intensity);
+    m_cn1->setArgs(m_scratchpads, m_states, m_intensity);
 
     m_cn2 = new Cn2RyoKernel(m_program);
-    m_cn2->setArgs(m_scratchpads, m_states, m_output, intensity);
+    m_cn2->setArgs(m_scratchpads, m_states, m_output, m_intensity);
 }
 
 
@@ -129,8 +124,6 @@ void xmrig::OclRyoRunner::init()
 {
     OclBaseRunner::init();
 
-    const size_t g_thd = data().thread.intensity();
-
-    m_scratchpads = createSubBuffer(CL_MEM_READ_WRITE, data().algorithm.l3() * g_thd);
-    m_states      = createSubBuffer(CL_MEM_READ_WRITE, 200 * g_thd);
+    m_scratchpads = createSubBuffer(CL_MEM_READ_WRITE, data().algorithm.l3() * m_intensity);
+    m_states      = createSubBuffer(CL_MEM_READ_WRITE, 200 * m_intensity);
 }

src/backend/opencl/runners/tools/OclCnR.cpp

@@ -237,7 +237,7 @@ private:
         for (size_t i = 0; i < OclCnR::kHeightChunkSize; ++i) {
             V4_Instruction code[256];
             const int code_size      = v4_random_math_init<Algorithm::CN_R>(code, offset + i);
-            const std::string kernel = std::regex_replace(cryptonight_r_cl, std::regex("XMRIG_INCLUDE_RANDOM_MATH"), getCode(code, code_size));
+            const std::string kernel = std::regex_replace(std::string(cryptonight_r_cl), std::regex("XMRIG_INCLUDE_RANDOM_MATH"), getCode(code, code_size));
 
             source += std::regex_replace(kernel, std::regex("KERNEL_NAME"), "cn1_" + std::to_string(offset + i));
         }

src/backend/opencl/runners/tools/OclSharedData.cpp

@@ -23,17 +23,44 @@
  */
 
 
-#include "backend/opencl/OclInterleave.h"
+#include "backend/opencl/runners/tools/OclSharedData.h"
+#include "backend/opencl/wrappers/OclLib.h"
 #include "base/io/log/Log.h"
 #include "base/tools/Chrono.h"
+#include "crypto/rx/Rx.h"
+#include "crypto/rx/RxDataset.h"
 
 
+#include <algorithm>
 #include <cinttypes>
+#include <stdexcept>
 #include <thread>
 
 
-uint64_t xmrig::OclInterleave::adjustDelay(size_t id)
+constexpr size_t oneGiB = 1024 * 1024 * 1024;
+
+
+cl_mem xmrig::OclSharedData::createBuffer(cl_context context, size_t size, size_t &offset)
 {
+    std::lock_guard<std::mutex> lock(m_mutex);
+
+    offset += size * m_offset++;
+    size   = std::max(size * m_threads, oneGiB);
+
+    if (!m_buffer) {
+        m_buffer = OclLib::createBuffer(context, CL_MEM_READ_WRITE, size);
+    }
+
+    return OclLib::retain(m_buffer);
+}
+
+
+uint64_t xmrig::OclSharedData::adjustDelay(size_t id)
+{
+    if (m_threads < 2) {
+        return 0;
+    }
+
     const uint64_t t0 = Chrono::steadyMSecs();
     uint64_t delay    = 0;
 
@@ -69,8 +96,12 @@ uint64_t xmrig::OclInterleave::adjustDelay(size_t id)
 }
 
 
-uint64_t xmrig::OclInterleave::resumeDelay(size_t id)
+uint64_t xmrig::OclSharedData::resumeDelay(size_t id)
 {
+    if (m_threads < 2) {
+        return 0;
+    }
+
     uint64_t delay = 0;
 
     {
@@ -99,14 +130,28 @@ uint64_t xmrig::OclInterleave::resumeDelay(size_t id)
 }
 
 
-void xmrig::OclInterleave::setResumeCounter(uint32_t value)
+void xmrig::OclSharedData::release()
 {
+    OclLib::release(m_buffer);
+
+#   ifdef XMRIG_ALGO_RANDOMX
+    OclLib::release(m_dataset);
+#   endif
+}
+
+
+void xmrig::OclSharedData::setResumeCounter(uint32_t value)
+{
+    if (m_threads < 2) {
+        return;
+    }
+
     std::lock_guard<std::mutex> lock(m_mutex);
     m_resumeCounter = value;
 }
 
 
-void xmrig::OclInterleave::setRunTime(uint64_t time)
+void xmrig::OclSharedData::setRunTime(uint64_t time)
 {
     // averagingBias = 1.0 - only the last delta time is taken into account
     // averagingBias = 0.5 - the last delta time has the same weight as all the previous ones combined
@@ -116,3 +161,34 @@ void xmrig::OclInterleave::setRunTime(uint64_t time)
     std::lock_guard<std::mutex> lock(m_mutex);
     m_averageRunTime = m_averageRunTime * (1.0 - averagingBias) + time * averagingBias;
 }
+
+
+#ifdef XMRIG_ALGO_RANDOMX
+cl_mem xmrig::OclSharedData::dataset() const
+{
+    if (!m_dataset) {
+        throw std::runtime_error("RandomX dataset is not available");
+    }
+
+    return OclLib::retain(m_dataset);
+}
+
+
+void xmrig::OclSharedData::createDataset(cl_context ctx, const Job &job, bool host)
+{
+    if (m_dataset) {
+        return;
+    }
+
+    cl_int ret;
+
+    if (host) {
+        auto dataset = Rx::dataset(job, 0);
+
+        m_dataset = OclLib::createBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, RxDataset::maxSize(), dataset->raw(), &ret);
+    }
+    else {
+        m_dataset = OclLib::createBuffer(ctx, CL_MEM_READ_ONLY, RxDataset::maxSize(), nullptr, &ret);
+    }
+}
+#endif

src/backend/opencl/runners/tools/OclSharedData.h

@@ -22,42 +22,62 @@
  *   along with this program. If not, see <http://www.gnu.org/licenses/>.
  */
 
-#ifndef XMRIG_OCLINTERLEAVE_H
-#define XMRIG_OCLINTERLEAVE_H
+#ifndef XMRIG_OCLSHAREDDATA_H
+#define XMRIG_OCLSHAREDDATA_H
 
 
 #include <memory>
 #include <mutex>
 
 
+using cl_context = struct _cl_context *;
+using cl_mem     = struct _cl_mem *;
+
+
 namespace xmrig {
 
 
-class OclInterleave
+class Job;
+
+
+class OclSharedData
 {
 public:
-    OclInterleave() = delete;
-    inline OclInterleave(size_t threads) : m_threads(threads) {}
+    OclSharedData() = default;
 
+    cl_mem createBuffer(cl_context context, size_t size, size_t &offset);
     uint64_t adjustDelay(size_t id);
     uint64_t resumeDelay(size_t id);
+    void release();
     void setResumeCounter(uint32_t value);
     void setRunTime(uint64_t time);
 
+    inline size_t threads() const       { return m_threads; }
+
+    inline OclSharedData &operator++()  { ++m_threads; return *this; }
+
+#   ifdef XMRIG_ALGO_RANDOMX
+    cl_mem dataset() const;
+    void createDataset(cl_context ctx, const Job &job, bool host);
+#   endif
+
 private:
-    const size_t m_threads;
+    cl_mem m_buffer           = nullptr;
     double m_averageRunTime   = 0.0;
     double m_threshold        = 0.95;
+    size_t m_offset           = 0;
+    size_t m_threads          = 0;
     std::mutex m_mutex;
     uint32_t m_resumeCounter  = 0;
     uint64_t m_timestamp      = 0;
+
+#   ifdef XMRIG_ALGO_RANDOMX
+    cl_mem m_dataset          = nullptr;
+#   endif
 };
 
 
-using OclInterleavePtr = std::shared_ptr<OclInterleave>;
-
-
 } /* namespace xmrig */
 
 
-#endif /* XMRIG_OCLINTERLEAVE_H */
+#endif /* XMRIG_OCLSHAREDDATA_H */

src/backend/opencl/runners/tools/OclSharedState.cpp

@@ -23,32 +23,52 @@
  */
 
 
-#include "backend/opencl/runners/tools/OclRxDataset.h"
-#include "backend/opencl/wrappers/OclLib.h"
-#include "crypto/rx/Rx.h"
-#include "crypto/rx/RxDataset.h"
+#include "backend/opencl/runners/tools/OclSharedState.h"
+#include "backend/opencl/runners/tools/OclSharedData.h"
 
 
-void xmrig::OclRxDataset::createBuffer(cl_context ctx, const Job &job, bool host)
+#include <cassert>
+#include <map>
+
+
+namespace xmrig {
+
+
+static std::map<uint32_t, OclSharedData> map;
+
+
+} // namespace xmrig
+
+
+xmrig::OclSharedData &xmrig::OclSharedState::get(uint32_t index)
 {
-    if (m_dataset) {
-        return;
-    }
-
-    cl_int ret;
-
-    if (host) {
-        auto dataset = Rx::dataset(job, 0);
-
-        m_dataset = OclLib::createBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, RxDataset::maxSize(), dataset->raw(), &ret);
-    }
-    else {
-        m_dataset = OclLib::createBuffer(ctx, CL_MEM_READ_ONLY, RxDataset::maxSize(), nullptr, &ret);
-    }
+    return map[index];
 }
 
 
-xmrig::OclRxDataset::~OclRxDataset()
+void xmrig::OclSharedState::release()
 {
-    OclLib::release(m_dataset);
+    for (auto &kv : map) {
+        kv.second.release();
+    }
+
+    map.clear();
+}
+
+
+void xmrig::OclSharedState::start(const std::vector<OclLaunchData> &threads, const Job &job)
+{
+    assert(map.empty());
+
+    for (const auto &data : threads) {
+        auto &sharedData = map[data.device.index()];
+
+        ++sharedData;
+
+#       ifdef XMRIG_ALGO_RANDOMX
+        if (data.algorithm.family() == Algorithm::RANDOM_X) {
+            sharedData.createDataset(data.ctx, job, data.thread.isDatasetHost());
+        }
+#       endif
+    }
 }

src/backend/opencl/runners/tools/OclSharedState.h

@@ -22,47 +22,26 @@
  *   along with this program. If not, see <http://www.gnu.org/licenses/>.
  */
 
-#ifndef XMRIG_OCLRXDATASET_H
-#define XMRIG_OCLRXDATASET_H
+#ifndef XMRIG_OCLSHAREDSTATE_H
+#define XMRIG_OCLSHAREDSTATE_H
 
 
-#include "base/tools/Object.h"
-
-
-#include <memory>
-
-
-using cl_context = struct _cl_context *;
-using cl_mem     = struct _cl_mem *;
+#include "backend/opencl/OclLaunchData.h"
 
 
 namespace xmrig {
 
 
-class Job;
-
-
-class OclRxDataset
+class OclSharedState
 {
 public:
-    XMRIG_DISABLE_COPY_MOVE(OclRxDataset)
-
-    OclRxDataset() = default;
-    ~OclRxDataset();
-
-    inline cl_mem get() const { return m_dataset; }
-
-    void createBuffer(cl_context ctx,const Job &job, bool host);
-
-private:
-    cl_mem m_dataset = nullptr;
+    static OclSharedData &get(uint32_t index);
+    static void release();
+    static void start(const std::vector<OclLaunchData> &threads, const Job &job);
 };
 
 
-using OclRxDatasetPtr = std::shared_ptr<OclRxDataset>;
-
-
 } /* namespace xmrig */
 
 
-#endif /* XMRIG_OCLINTERLEAVE_H */
+#endif /* XMRIG_OCLSHAREDSTATE_H */

src/backend/opencl/wrappers/OclContext.cpp

@@ -24,6 +24,7 @@
 
 
 #include "backend/opencl/wrappers/OclContext.h"
+#include "backend/opencl/runners/tools/OclSharedState.h"
 #include "backend/opencl/wrappers/OclLib.h"
 
 
@@ -42,7 +43,7 @@ xmrig::OclContext::~OclContext()
 }
 
 
-bool xmrig::OclContext::init(const std::vector<OclDevice> &devices, std::vector<OclLaunchData> &threads, const Job &job)
+bool xmrig::OclContext::init(const std::vector<OclDevice> &devices, std::vector<OclLaunchData> &threads)
 {
     if (!m_ctx) {
         std::vector<cl_device_id> ids(devices.size());
@@ -59,12 +60,6 @@ bool xmrig::OclContext::init(const std::vector<OclDevice> &devices, std::vector<
 
     for (OclLaunchData &data : threads) {
         data.ctx = m_ctx;
-
-#       ifdef XMRIG_ALGO_RANDOMX
-        if (data.algorithm.family() == Algorithm::RANDOM_X) {
-            data.dataset->createBuffer(m_ctx, job, data.thread.isDatasetHost());
-        }
-#       endif
     }
 
     return true;

src/backend/opencl/wrappers/OclContext.h

@@ -49,7 +49,7 @@ public:
     OclContext(const OclDevice &device);
     ~OclContext();
 
-    bool init(const std::vector<OclDevice> &devices, std::vector<OclLaunchData> &threads, const Job &job);
+    bool init(const std::vector<OclDevice> &devices, std::vector<OclLaunchData> &threads);
 
     inline bool isValid() const     { return m_ctx != nullptr; }
     inline cl_context ctx() const   { return m_ctx; }

src/backend/opencl/wrappers/OclDevice.cpp

@@ -59,7 +59,7 @@ extern bool ocl_vega_cn_generator(const OclDevice &device, const Algorithm &algo
 extern bool ocl_generic_cn_generator(const OclDevice &device, const Algorithm &algorithm, OclThreads &threads);
 
 
-ocl_gen_config_fun generators[] = {
+static ocl_gen_config_fun generators[] = {
 #   ifdef XMRIG_ALGO_RANDOMX
     ocl_generic_rx_generator,
 #   endif
@@ -142,16 +142,14 @@ xmrig::OclDevice::OclDevice(uint32_t index, cl_device_id id, cl_platform_id plat
         topology_amd topology;
 
         if (OclLib::getDeviceInfo(id, 0x4037 /* CL_DEVICE_TOPOLOGY_AMD */, sizeof(topology), &topology, nullptr) == CL_SUCCESS && topology.raw.type == 1) {
-            m_topology    = true;
-            m_pciTopology = PciTopology(static_cast<uint32_t>(topology.pcie.bus), static_cast<uint32_t>(topology.pcie.device), static_cast<uint32_t>(topology.pcie.function));
+            m_topology = PciTopology(static_cast<uint32_t>(topology.pcie.bus), static_cast<uint32_t>(topology.pcie.device), static_cast<uint32_t>(topology.pcie.function));
         }
     }
     else if (m_vendorId == OCL_VENDOR_NVIDIA) {
         cl_uint bus = 0;
         if (OclLib::getDeviceInfo(id, 0x4008 /* CL_DEVICE_PCI_BUS_ID_NV */, sizeof (bus), &bus, nullptr) == CL_SUCCESS) {
-            m_topology    = true;
             cl_uint slot  = OclLib::getUint(id, 0x4009 /* CL_DEVICE_PCI_SLOT_ID_NV */);
-            m_pciTopology = PciTopology(bus, (slot >> 3) & 0xff, slot & 7);
+            m_topology = PciTopology(bus, (slot >> 3) & 0xff, slot & 7);
         }
     }
 }
@@ -205,3 +203,18 @@ void xmrig::OclDevice::generate(const Algorithm &algorithm, OclThreads &threads)
         }
     }
 }
+
+
+#ifdef XMRIG_FEATURE_API
+void xmrig::OclDevice::toJSON(rapidjson::Value &out, rapidjson::Document &doc) const
+{
+    using namespace rapidjson;
+    auto &allocator = doc.GetAllocator();
+
+    out.AddMember("board",       board().toJSON(doc), allocator);
+    out.AddMember("name",        name().toJSON(doc), allocator);
+    out.AddMember("bus_id",      topology().toString().toJSON(doc), allocator);
+    out.AddMember("cu",          computeUnits(), allocator);
+    out.AddMember("global_mem",  static_cast<uint64_t>(globalMemSize()), allocator);
+}
+#endif

src/backend/opencl/wrappers/OclDevice.h

@@ -69,10 +69,9 @@ public:
     uint32_t clock() const;
     void generate(const Algorithm &algorithm, OclThreads &threads) const;
 
-    inline bool hasTopology() const             { return m_topology; }
     inline bool isValid() const                 { return m_id != nullptr && m_platform != nullptr; }
     inline cl_device_id id() const              { return m_id; }
-    inline const PciTopology &topology() const  { return m_pciTopology; }
+    inline const PciTopology &topology() const  { return m_topology; }
     inline const String &board() const          { return m_board.isNull() ? m_name : m_board; }
     inline const String &name() const           { return m_name; }
     inline const String &vendor() const         { return m_vendor; }
@@ -81,8 +80,11 @@ public:
     inline uint32_t computeUnits() const        { return m_computeUnits; }
     inline uint32_t index() const               { return m_index; }
 
+#   ifdef XMRIG_FEATURE_API
+    void toJSON(rapidjson::Value &out, rapidjson::Document &doc) const;
+#   endif
+
 private:
-    bool m_topology                 = false;
     cl_device_id m_id               = nullptr;
     cl_platform_id m_platform       = nullptr;
     const String m_board;
@@ -91,7 +93,7 @@ private:
     const uint32_t m_computeUnits   = 1;
     const uint32_t m_index          = 0;
     OclVendor m_vendorId            = OCL_VENDOR_UNKNOWN;
-    PciTopology m_pciTopology;
+    PciTopology m_topology;
     Type m_type                     = Unknown;
 };
 

src/backend/opencl/wrappers/OclLib.cpp

@@ -73,6 +73,8 @@ static const char *kReleaseDevice                    = "clReleaseDevice";
 static const char *kReleaseKernel                    = "clReleaseKernel";
 static const char *kReleaseMemObject                 = "clReleaseMemObject";
 static const char *kReleaseProgram                   = "clReleaseProgram";
+static const char *kRetainMemObject                  = "clRetainMemObject";
+static const char *kRetainProgram                    = "clRetainProgram";
 static const char *kSetKernelArg                     = "clSetKernelArg";
 static const char *kSetMemObjectDestructorCallback   = "clSetMemObjectDestructorCallback";
 static const char *kUnloadPlatformCompiler           = "clUnloadPlatformCompiler";
@@ -105,6 +107,8 @@ typedef cl_int (CL_API_CALL *releaseDevice_t)(cl_device_id device);
 typedef cl_int (CL_API_CALL *releaseKernel_t)(cl_kernel);
 typedef cl_int (CL_API_CALL *releaseMemObject_t)(cl_mem);
 typedef cl_int (CL_API_CALL *releaseProgram_t)(cl_program);
+typedef cl_int (CL_API_CALL *retainMemObject_t)(cl_mem);
+typedef cl_int (CL_API_CALL *retainProgram_t)(cl_program);
 typedef cl_int (CL_API_CALL *setKernelArg_t)(cl_kernel, cl_uint, size_t, const void *);
 typedef cl_int (CL_API_CALL *setMemObjectDestructorCallback_t)(cl_mem, void (CL_CALLBACK *)(cl_mem, void *), void *);
 typedef cl_int (CL_API_CALL *unloadPlatformCompiler_t)(cl_platform_id);
@@ -146,6 +150,8 @@ static releaseDevice_t pReleaseDevice                                       = nu
 static releaseKernel_t pReleaseKernel                                       = nullptr;
 static releaseMemObject_t pReleaseMemObject                                 = nullptr;
 static releaseProgram_t pReleaseProgram                                     = nullptr;
+static retainMemObject_t pRetainMemObject                                   = nullptr;
+static retainProgram_t pRetainProgram                                       = nullptr;
 static setKernelArg_t pSetKernelArg                                         = nullptr;
 static setMemObjectDestructorCallback_t pSetMemObjectDestructorCallback     = nullptr;
 static unloadPlatformCompiler_t pUnloadPlatformCompiler                     = nullptr;
@@ -235,6 +241,8 @@ bool xmrig::OclLib::load()
     DLSYM(UnloadPlatformCompiler);
     DLSYM(SetMemObjectDestructorCallback);
     DLSYM(CreateSubBuffer);
+    DLSYM(RetainProgram);
+    DLSYM(RetainMemObject);
 
 #   if defined(CL_VERSION_2_0)
     uv_dlsym(&oclLib, kCreateCommandQueueWithProperties, reinterpret_cast<void**>(&pCreateCommandQueueWithProperties));
@@ -666,6 +674,18 @@ cl_mem xmrig::OclLib::createSubBuffer(cl_mem buffer, cl_mem_flags flags, size_t
 }
 
 
+cl_mem xmrig::OclLib::retain(cl_mem memobj) noexcept
+{
+    assert(pRetainMemObject != nullptr);
+
+    if (memobj != nullptr) {
+        pRetainMemObject(memobj);
+    }
+
+    return memobj;
+}
+
+
 cl_program xmrig::OclLib::createProgramWithBinary(cl_context context, cl_uint num_devices, const cl_device_id *device_list, const size_t *lengths, const unsigned char **binaries, cl_int *binary_status, cl_int *errcode_ret) noexcept
 {
     assert(pCreateProgramWithBinary != nullptr);
@@ -696,6 +716,18 @@ cl_program xmrig::OclLib::createProgramWithSource(cl_context context, cl_uint co
 }
 
 
+cl_program xmrig::OclLib::retain(cl_program program) noexcept
+{
+    assert(pRetainProgram != nullptr);
+
+    if (program != nullptr) {
+        pRetainProgram(program);
+    }
+
+    return program;
+}
+
+
 cl_uint xmrig::OclLib::getNumPlatforms() noexcept
 {
     cl_uint count = 0;

src/backend/opencl/wrappers/OclLib.h

@@ -79,8 +79,10 @@ public:
     static cl_mem createBuffer(cl_context context, cl_mem_flags flags, size_t size, void *host_ptr, cl_int *errcode_ret) noexcept;
     static cl_mem createSubBuffer(cl_mem buffer, cl_mem_flags flags, size_t offset, size_t size, cl_int *errcode_ret) noexcept;
     static cl_mem createSubBuffer(cl_mem buffer, cl_mem_flags flags, size_t offset, size_t size);
+    static cl_mem retain(cl_mem memobj) noexcept;
     static cl_program createProgramWithBinary(cl_context context, cl_uint num_devices, const cl_device_id *device_list, const size_t *lengths, const unsigned char **binaries, cl_int *binary_status, cl_int *errcode_ret) noexcept;
     static cl_program createProgramWithSource(cl_context context, cl_uint count, const char **strings, const size_t *lengths, cl_int *errcode_ret) noexcept;
+    static cl_program retain(cl_program program) noexcept;
     static cl_uint getNumPlatforms() noexcept;
     static cl_uint getUint(cl_command_queue command_queue, cl_command_queue_info param_name, cl_uint defaultValue = 0) noexcept;
     static cl_uint getUint(cl_context context, cl_context_info param_name, cl_uint defaultValue = 0) noexcept;

src/backend/opencl/wrappers/OclPlatform.cpp

@@ -73,7 +73,7 @@ rapidjson::Value xmrig::OclPlatform::toJSON(rapidjson::Document &doc) const
     }
 
     Value out(kObjectType);
-    out.AddMember("index",      index(), allocator);
+    out.AddMember("index",      static_cast<uint64_t>(index()), allocator);
     out.AddMember("profile",    profile().toJSON(doc), allocator);
     out.AddMember("version",    version().toJSON(doc), allocator);
     out.AddMember("name",       name().toJSON(doc), allocator);

src/base/api/Api.cpp

@@ -117,9 +117,10 @@ void xmrig::Api::exec(IApiRequest &request)
         auto &allocator = request.doc().GetAllocator();
 
         request.accept();
-        request.reply().AddMember("id",        StringRef(m_id),       allocator);
-        request.reply().AddMember("worker_id", StringRef(m_workerId), allocator);
-        request.reply().AddMember("uptime",    (Chrono::currentMSecsSinceEpoch() - m_timestamp) / 1000, allocator);
+        request.reply().AddMember("id",         StringRef(m_id),       allocator);
+        request.reply().AddMember("worker_id",  StringRef(m_workerId), allocator);
+        request.reply().AddMember("uptime",     (Chrono::currentMSecsSinceEpoch() - m_timestamp) / 1000, allocator);
+        request.reply().AddMember("restricted", request.isRestricted(), allocator);
 
         Value features(kArrayType);
 #       ifdef XMRIG_FEATURE_API

src/base/api/Api.h

@@ -27,10 +27,11 @@
 
 
 #include <vector>
-#include <stdint.h>
+#include <cstdint>
 
 
 #include "base/kernel/interfaces/IBaseListener.h"
+#include "base/tools/Object.h"
 
 
 namespace xmrig {
@@ -47,6 +48,8 @@ class String;
 class Api : public IBaseListener
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(Api)
+
     Api(Base *base);
     ~Api() override;
 

src/base/io/Console.cpp

@@ -31,8 +31,11 @@
 xmrig::Console::Console(IConsoleListener *listener)
     : m_listener(listener)
 {
-    m_tty = new uv_tty_t;
+    if (!isSupported()) {
+        return;
+    }
 
+    m_tty = new uv_tty_t;
     m_tty->data = this;
     uv_tty_init(uv_default_loop(), m_tty, 0, 1);
 
@@ -53,6 +56,10 @@ xmrig::Console::~Console()
 
 void xmrig::Console::stop()
 {
+    if (!m_tty) {
+        return;
+    }
+
     uv_tty_reset_mode();
 
     Handle::close(m_tty);
@@ -60,6 +67,13 @@ void xmrig::Console::stop()
 }
 
 
+bool xmrig::Console::isSupported() const
+{
+    const uv_handle_type type = uv_guess_handle(0);
+    return type == UV_TTY || type == UV_NAMED_PIPE;
+}
+
+
 void xmrig::Console::onAllocBuffer(uv_handle_t *handle, size_t, uv_buf_t *buf)
 {
     auto console = static_cast<Console*>(handle->data);

src/base/io/Console.h

@@ -26,9 +26,11 @@
 #define XMRIG_CONSOLE_H
 
 
-#include <uv.h>
+#include "base/tools/Object.h"
 
 
+#include <uv.h>
+
 
 namespace xmrig {
 
@@ -39,18 +41,22 @@ class IConsoleListener;
 class Console
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(Console)
+
     Console(IConsoleListener *listener);
     ~Console();
 
     void stop();
 
 private:
+    bool isSupported() const;
+
     static void onAllocBuffer(uv_handle_t *handle, size_t suggested_size, uv_buf_t *buf);
     static void onRead(uv_stream_t *stream, ssize_t nread, const uv_buf_t *buf);
 
-    char m_buf[1];
+    char m_buf[1] = { 0 };
     IConsoleListener *m_listener;
-    uv_tty_t *m_tty;
+    uv_tty_t *m_tty = nullptr;
 };
 
 

src/base/io/json/Json.cpp

@@ -27,6 +27,9 @@
 #include "rapidjson/document.h"
 
 
+#include <cassert>
+
+
 namespace xmrig {
 
 static const rapidjson::Value kNullValue;
@@ -36,6 +39,8 @@ static const rapidjson::Value kNullValue;
 
 bool xmrig::Json::getBool(const rapidjson::Value &obj, const char *key, bool defaultValue)
 {
+    assert(obj.IsObject());
+
     auto i = obj.FindMember(key);
     if (i != obj.MemberEnd() && i->value.IsBool()) {
         return i->value.GetBool();
@@ -47,6 +52,8 @@ bool xmrig::Json::getBool(const rapidjson::Value &obj, const char *key, bool def
 
 const char *xmrig::Json::getString(const rapidjson::Value &obj, const char *key,  const char *defaultValue)
 {
+    assert(obj.IsObject());
+
     auto i = obj.FindMember(key);
     if (i != obj.MemberEnd() && i->value.IsString()) {
         return i->value.GetString();
@@ -58,6 +65,8 @@ const char *xmrig::Json::getString(const rapidjson::Value &obj, const char *key,
 
 const rapidjson::Value &xmrig::Json::getArray(const rapidjson::Value &obj, const char *key)
 {
+    assert(obj.IsObject());
+
     auto i = obj.FindMember(key);
     if (i != obj.MemberEnd() && i->value.IsArray()) {
         return i->value;
@@ -69,6 +78,8 @@ const rapidjson::Value &xmrig::Json::getArray(const rapidjson::Value &obj, const
 
 const rapidjson::Value &xmrig::Json::getObject(const rapidjson::Value &obj, const char *key)
 {
+    assert(obj.IsObject());
+
     auto i = obj.FindMember(key);
     if (i != obj.MemberEnd() && i->value.IsObject()) {
         return i->value;
@@ -80,6 +91,8 @@ const rapidjson::Value &xmrig::Json::getObject(const rapidjson::Value &obj, cons
 
 const rapidjson::Value &xmrig::Json::getValue(const rapidjson::Value &obj, const char *key)
 {
+    assert(obj.IsObject());
+
     auto i = obj.FindMember(key);
     if (i != obj.MemberEnd()) {
         return i->value;
@@ -91,6 +104,8 @@ const rapidjson::Value &xmrig::Json::getValue(const rapidjson::Value &obj, const
 
 int xmrig::Json::getInt(const rapidjson::Value &obj, const char *key, int defaultValue)
 {
+    assert(obj.IsObject());
+
     auto i = obj.FindMember(key);
     if (i != obj.MemberEnd() && i->value.IsInt()) {
         return i->value.GetInt();
@@ -102,6 +117,8 @@ int xmrig::Json::getInt(const rapidjson::Value &obj, const char *key, int defaul
 
 int64_t xmrig::Json::getInt64(const rapidjson::Value &obj, const char *key, int64_t defaultValue)
 {
+    assert(obj.IsObject());
+
     auto i = obj.FindMember(key);
     if (i != obj.MemberEnd() && i->value.IsInt64()) {
         return i->value.GetInt64();
@@ -113,6 +130,8 @@ int64_t xmrig::Json::getInt64(const rapidjson::Value &obj, const char *key, int6
 
 uint64_t xmrig::Json::getUint64(const rapidjson::Value &obj, const char *key, uint64_t defaultValue)
 {
+    assert(obj.IsObject());
+
     auto i = obj.FindMember(key);
     if (i != obj.MemberEnd() && i->value.IsUint64()) {
         return i->value.GetUint64();
@@ -124,6 +143,8 @@ uint64_t xmrig::Json::getUint64(const rapidjson::Value &obj, const char *key, ui
 
 unsigned xmrig::Json::getUint(const rapidjson::Value &obj, const char *key, unsigned defaultValue)
 {
+    assert(obj.IsObject());
+
     auto i = obj.FindMember(key);
     if (i != obj.MemberEnd() && i->value.IsUint()) {
         return i->value.GetUint();

src/base/io/json/Json_win.cpp

@@ -28,6 +28,7 @@
 
 #ifdef __GNUC__
 #   include <fcntl.h>
+#   include <sys/stat.h>
 #   include <ext/stdio_filebuf.h>
 #endif
 
@@ -102,7 +103,7 @@ bool xmrig::Json::save(const char *fileName, const rapidjson::Document &doc)
         return false;
     }
 #   elif defined(__GNUC__)
-    const int fd = _wopen(toUtf16(fileName).c_str(), _O_WRONLY | _O_BINARY | _O_CREAT | _O_TRUNC);
+    const int fd = _wopen(toUtf16(fileName).c_str(), _O_WRONLY | _O_BINARY | _O_CREAT | _O_TRUNC, _S_IWRITE);
     if (fd == -1) {
         return false;
     }

src/base/io/log/backends/ConsoleLog.cpp

@@ -24,7 +24,7 @@
  */
 
 
-#include <stdio.h>
+#include <cstdio>
 
 
 #include "base/tools/Handle.h"
@@ -32,9 +32,13 @@
 #include "base/io/log/Log.h"
 
 
-xmrig::ConsoleLog::ConsoleLog() :
-    m_stream(nullptr)
+xmrig::ConsoleLog::ConsoleLog()
 {
+    if (!isSupported()) {
+        Log::colors = false;
+        return;
+    }
+
     m_tty = new uv_tty_t;
 
     if (uv_tty_init(uv_default_loop(), m_tty, 1, 0) < 0) {
@@ -66,7 +70,7 @@ xmrig::ConsoleLog::~ConsoleLog()
 
 void xmrig::ConsoleLog::print(int, const char *line, size_t, size_t size, bool colors)
 {
-    if (Log::colors != colors) {
+    if (!m_tty || Log::colors != colors) {
         return;
     }
 
@@ -86,12 +90,18 @@ void xmrig::ConsoleLog::print(int, const char *line, size_t, size_t size, bool c
 }
 
 
+bool xmrig::ConsoleLog::isSupported() const
+{
+    const uv_handle_type type = uv_guess_handle(1);
+    return type == UV_TTY || type == UV_NAMED_PIPE;
+}
+
+
 bool xmrig::ConsoleLog::isWritable() const
 {
     if (!m_stream || uv_is_writable(m_stream) != 1) {
         return false;
     }
 
-    const uv_handle_type type = uv_guess_handle(1);
-    return type == UV_TTY || type == UV_NAMED_PIPE;
+    return isSupported();
 }

src/base/io/log/backends/ConsoleLog.h

@@ -27,11 +27,12 @@
 #define XMRIG_CONSOLELOG_H
 
 
-typedef struct uv_stream_s uv_stream_t;
-typedef struct uv_tty_s uv_tty_t;
+using uv_stream_t = struct uv_stream_s;
+using uv_tty_t    = struct uv_tty_s;
 
 
 #include "base/kernel/interfaces/ILogBackend.h"
+#include "base/tools/Object.h"
 
 
 namespace xmrig {
@@ -40,6 +41,8 @@ namespace xmrig {
 class ConsoleLog : public ILogBackend
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE(ConsoleLog)
+
     ConsoleLog();
     ~ConsoleLog() override;
 
@@ -47,10 +50,11 @@ protected:
     void print(int level, const char *line, size_t offset, size_t size, bool colors) override;
 
 private:
+    bool isSupported() const;
     bool isWritable() const;
 
-    uv_stream_t *m_stream;
-    uv_tty_t *m_tty;
+    uv_stream_t *m_stream = nullptr;
+    uv_tty_t *m_tty       = nullptr;
 };
 
 

src/base/kernel/Base.cpp

@@ -23,7 +23,7 @@
  */
 
 
-#include <assert.h>
+#include <cassert>
 #include <memory>
 
 
@@ -64,15 +64,16 @@ static const char *kConfigPathV2 = "/2/config";
 #endif
 
 
-class xmrig::BasePrivate
+namespace xmrig {
+
+
+class BasePrivate
 {
 public:
-    inline BasePrivate(Process *process) :
-        api(nullptr),
-        config(nullptr),
-        process(process),
-        watcher(nullptr)
-    {}
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(BasePrivate)
+
+
+    inline BasePrivate(Process *process) : config(load(process)) {}
 
 
     inline ~BasePrivate()
@@ -94,13 +95,33 @@ public:
     }
 
 
-    inline Config *load()
+    inline void replace(Config *newConfig)
+    {
+        Config *previousConfig = config;
+        config = newConfig;
+
+        for (IBaseListener *listener : listeners) {
+            listener->onConfigChanged(config, previousConfig);
+        }
+
+        delete previousConfig;
+    }
+
+
+    Api *api            = nullptr;
+    Config *config      = nullptr;
+    std::vector<IBaseListener *> listeners;
+    Watcher *watcher    = nullptr;
+
+
+private:
+    inline Config *load(Process *process)
     {
         JsonChain chain;
         ConfigTransform transform;
         std::unique_ptr<Config> config;
 
-        transform.load(chain, process, transform);
+        ConfigTransform::load(chain, process, transform);
 
         if (read(chain, config)) {
             return config.release();
@@ -122,29 +143,12 @@ public:
 
         return nullptr;
     }
-
-
-    inline void replace(Config *newConfig)
-    {
-        Config *previousConfig = config;
-        config = newConfig;
-
-        for (IBaseListener *listener : listeners) {
-            listener->onConfigChanged(config, previousConfig);
-        }
-
-        delete previousConfig;
-    }
-
-
-    Api *api;
-    Config *config;
-    Process *process;
-    std::vector<IBaseListener *> listeners;
-    Watcher *watcher;
 };
 
 
+} // namespace xmrig
+
+
 xmrig::Base::Base(Process *process)
     : d_ptr(new BasePrivate(process))
 {
@@ -165,14 +169,6 @@ bool xmrig::Base::isReady() const
 
 int xmrig::Base::init()
 {
-    d_ptr->config = d_ptr->load();
-
-    if (!d_ptr->config) {
-        LOG_EMERG("No valid configuration found. Exiting.");
-
-        return 1;
-    }
-
 #   ifdef XMRIG_FEATURE_API
     d_ptr->api = new Api(this);
     d_ptr->api->addListener(this);
@@ -184,7 +180,7 @@ int xmrig::Base::init()
     Platform::setProcessPriority(config()->cpu().priority());
 #   endif
 
-    if (config()->isBackground()) {
+    if (isBackground()) {
         Log::background = true;
     }
     else {
@@ -240,6 +236,12 @@ xmrig::Api *xmrig::Base::api() const
 }
 
 
+bool xmrig::Base::isBackground() const
+{
+    return d_ptr->config && d_ptr->config->isBackground();
+}
+
+
 bool xmrig::Base::reload(const rapidjson::Value &json)
 {
     JsonReader reader(json);
@@ -247,7 +249,7 @@ bool xmrig::Base::reload(const rapidjson::Value &json)
         return false;
     }
 
-    Config *config = new Config();
+    auto config = new Config();
     if (!config->read(reader, d_ptr->config->fileName())) {
         delete config;
 
@@ -289,7 +291,7 @@ void xmrig::Base::onFileChanged(const String &fileName)
     JsonChain chain;
     chain.addFile(fileName);
 
-    Config *config = new Config();
+    auto config = new Config();
 
     if (!config->read(chain, chain.fileName())) {
         LOG_ERR("reloading failed");

src/base/kernel/Base.h

@@ -29,6 +29,7 @@
 #include "base/api/interfaces/IApiListener.h"
 #include "base/kernel/interfaces/IConfigListener.h"
 #include "base/kernel/interfaces/IWatcherListener.h"
+#include "base/tools/Object.h"
 #include "rapidjson/fwd.h"
 
 
@@ -45,6 +46,8 @@ class Process;
 class Base : public IWatcherListener, public IApiListener
 {
 public:
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(Base)
+
     Base(Process *process);
     ~Base() override;
 
@@ -54,6 +57,7 @@ public:
     virtual void stop();
 
     Api *api() const;
+    bool isBackground() const;
     bool reload(const rapidjson::Value &json);
     Config *config() const;
     void addListener(IBaseListener *listener);

src/base/kernel/Entry.cpp

@@ -23,7 +23,7 @@
  */
 
 
-#include <stdio.h>
+#include <cstdio>
 #include <uv.h>
 
 
@@ -161,7 +161,7 @@ int xmrig::Entry::exec(const Process &process, Id id)
 {
     switch (id) {
     case Usage:
-        printf(usage);
+        printf("%s\n", usage().c_str());
         return 0;
 
     case Version:

src/base/kernel/Process.cpp

@@ -23,8 +23,8 @@
  */
 
 
+#include <ctime>
 #include <uv.h>
-#include <time.h>
 
 
 #include "base/kernel/Process.h"
@@ -55,11 +55,6 @@ xmrig::Process::Process(int argc, char **argv) :
 }
 
 
-xmrig::Process::~Process()
-{
-}
-
-
 xmrig::String xmrig::Process::location(Location location, const char *fileName) const
 {
     constexpr const size_t max = 520;

src/base/kernel/Process.h

@@ -47,7 +47,6 @@ public:
 #   endif
 
     Process(int argc, char **argv);
-    ~Process();
 
     String location(Location location, const char *fileName = nullptr) const;
 

src/base/kernel/config/BaseTransform.cpp

@@ -23,7 +23,7 @@
  */
 
 
-#include <stdio.h>
+#include <cstdio>
 
 
 #ifdef _MSC_VER
@@ -47,15 +47,11 @@ namespace xmrig
 
 static const char *kAlgo  = "algo";
 static const char *kApi   = "api";
+static const char *kCoin  = "coin";
 static const char *kHttp  = "http";
 static const char *kPools = "pools";
 
-}
-
-
-xmrig::BaseTransform::BaseTransform()
-{
-}
+} // namespace xmrig
 
 
 void xmrig::BaseTransform::load(JsonChain &chain, Process *process, IConfigTransform &transform)
@@ -68,7 +64,7 @@ void xmrig::BaseTransform::load(JsonChain &chain, Process *process, IConfigTrans
 
     Document doc(kObjectType);
 
-    while (1) {
+    while (true) {
         key = getopt_long(argc, argv, short_options, options, nullptr);
         if (key < 0) {
             break;
@@ -107,6 +103,19 @@ void xmrig::BaseTransform::finalize(rapidjson::Document &doc)
             }
         }
     }
+
+    if (m_coin.isValid() && doc.HasMember(kPools)) {
+        auto &pools = doc[kPools];
+        for (Value &pool : pools.GetArray()) {
+            if (!pool.HasMember(kCoin)) {
+                pool.AddMember(StringRef(kCoin), m_coin.toJSON(), allocator);
+            }
+        }
+    }
+
+    if (m_http) {
+        set(doc, kHttp, "enabled", true);
+    }
 }
 
 
@@ -122,6 +131,15 @@ void xmrig::BaseTransform::transform(rapidjson::Document &doc, int key, const ch
         }
         break;
 
+    case IConfig::CoinKey: /* --coin */
+        if (!doc.HasMember(kPools)) {
+            m_coin = arg;
+        }
+        else {
+            return add(doc, kPools, kCoin, arg);
+        }
+        break;
+
     case IConfig::UserpassKey: /* --userpass */
         {
             const char *p = strrchr(arg, ':');
@@ -169,9 +187,11 @@ void xmrig::BaseTransform::transform(rapidjson::Document &doc, int key, const ch
         return set(doc, "log-file", arg);
 
     case IConfig::HttpAccessTokenKey: /* --http-access-token */
+        m_http = true;
         return set(doc, kHttp, "access-token", arg);
 
     case IConfig::HttpHostKey: /* --http-host */
+        m_http = true;
         return set(doc, kHttp, "host", arg);
 
     case IConfig::ApiWorkerIdKey: /* --api-worker-id */
@@ -226,8 +246,10 @@ void xmrig::BaseTransform::transformBoolean(rapidjson::Document &doc, int key, b
     case IConfig::TlsKey: /* --tls */
         return add(doc, kPools, "tls", enable);
 
+#   ifdef XMRIG_FEATURE_HTTP
     case IConfig::DaemonKey: /* --daemon */
         return add(doc, kPools, "daemon", enable);
+#   endif
 
 #   ifndef XMRIG_PROXY_PROJECT
     case IConfig::NicehashKey: /* --nicehash */
@@ -238,10 +260,12 @@ void xmrig::BaseTransform::transformBoolean(rapidjson::Document &doc, int key, b
         return set(doc, "colors", enable);
 
     case IConfig::HttpRestrictedKey: /* --http-no-restricted */
+        m_http = true;
         return set(doc, kHttp, "restricted", enable);
 
     case IConfig::HttpEnabledKey: /* --http-enabled */
-        return set(doc, kHttp, "enabled", enable);
+        m_http = true;
+        break;
 
     case IConfig::DryRunKey: /* --dry-run */
         return set(doc, "dry-run", enable);
@@ -268,13 +292,16 @@ void xmrig::BaseTransform::transformUint64(rapidjson::Document &doc, int key, ui
         return set(doc, "donate-over-proxy", arg);
 
     case IConfig::HttpPort: /* --http-port */
+        m_http = true;
         return set(doc, kHttp, "port", arg);
 
     case IConfig::PrintTimeKey: /* --print-time */
         return set(doc, "print-time", arg);
 
+#   ifdef XMRIG_FEATURE_HTTP
     case IConfig::DaemonPollKey:  /* --daemon-poll-interval */
         return add(doc, kPools, "daemon-poll-interval", arg);
+#   endif
 
     default:
         break;

src/base/kernel/config/BaseTransform.h

@@ -27,6 +27,7 @@
 
 
 #include "base/kernel/interfaces/IConfigTransform.h"
+#include "crypto/common/Coin.h"
 #include "rapidjson/document.h"
 
 
@@ -44,8 +45,6 @@ class Process;
 class BaseTransform : public IConfigTransform
 {
 public:
-    BaseTransform();
-
     static void load(JsonChain &chain, Process *process, IConfigTransform &transform);
 
 protected:
@@ -99,11 +98,14 @@ protected:
 
 protected:
     Algorithm m_algorithm;
+    Coin m_coin;
 
 
 private:
     void transformBoolean(rapidjson::Document &doc, int key, bool enable);
     void transformUint64(rapidjson::Document &doc, int key, uint64_t arg);
+
+    bool m_http = false;
 };
 
 

src/base/kernel/interfaces/IConfig.h

@@ -43,6 +43,7 @@ public:
     enum Keys {
         // common
         AlgorithmKey         = 'a',
+        CoinKey              = 1025,
         ApiWorkerIdKey       = 4002,
         ApiIdKey             = 4005,
         HttpPort             = 4100,
@@ -78,6 +79,7 @@ public:
         PrintTimeKey         = 1007,
 
         // xmrig cpu
+        CPUKey               = 1024,
         AVKey                = 'v',
         CPUAffinityKey       = 1020,
         DryRunKey            = 5000,
@@ -86,6 +88,7 @@ public:
         AssemblyKey          = 1015,
         RandomXInitKey       = 1022,
         RandomXNumaKey       = 1023,
+        CPUMaxThreadsKey     = 1026,
 
         // xmrig amd
         OclPlatformKey       = 1400,
@@ -99,6 +102,7 @@ public:
         OclMemChunkKey       = 1408,
         OclUnrollKey         = 1409,
         OclCompModeKey       = 1410,
+        OclKey               = 1411,
 
         // xmrig-proxy
         AccessLogFileKey     = 'A',

src/base/net/stratum/Client.cpp

@@ -334,6 +334,9 @@ bool xmrig::Client::parseJob(const rapidjson::Value &params, int *code)
     if (algo) {
         job.setAlgorithm(algo);
     }
+    else if (m_pool.coin().isValid()) {
+        job.setAlgorithm(m_pool.coin().algorithm(job.blob()[0]));
+    }
 
     job.setHeight(Json::getUint64(params, "height"));
 
@@ -426,7 +429,12 @@ bool xmrig::Client::verifyAlgorithm(const Algorithm &algorithm, const char *algo
 {
     if (!algorithm.isValid()) {
         if (!isQuiet()) {
-            LOG_ERR("[%s] Unknown/unsupported algorithm \"%s\" detected, reconnect", url(), algo);
+            if (algo == nullptr) {
+                LOG_ERR("[%s] unknown algorithm, make sure you set \"algo\" or \"coin\" option", url(), algo);
+            }
+            else {
+                LOG_ERR("[%s] unsupported algorithm \"%s\" detected, reconnect", url(), algo);
+            }
         }
 
         return false;
@@ -436,7 +444,7 @@ bool xmrig::Client::verifyAlgorithm(const Algorithm &algorithm, const char *algo
     m_listener->onVerifyAlgorithm(this, algorithm, &ok);
 
     if (!ok && !isQuiet()) {
-        LOG_ERR("[%s] Incompatible/disabled algorithm \"%s\" detected, reconnect", url(), algorithm.shortName());
+        LOG_ERR("[%s] incompatible/disabled algorithm \"%s\" detected, reconnect", url(), algorithm.shortName());
     }
 
     return ok;

src/base/net/stratum/DaemonClient.cpp

@@ -25,7 +25,7 @@
 
 
 #include <algorithm>
-#include <assert.h>
+#include <cassert>
 
 
 #include "3rdparty/http-parser/http_parser.h"
@@ -225,6 +225,10 @@ bool xmrig::DaemonClient::parseJob(const rapidjson::Value &params, int *code)
     job.setDiff(Json::getUint64(params, "difficulty"));
     job.setId(blocktemplate.data() + blocktemplate.size() - 32);
 
+    if (m_pool.coin().isValid()) {
+        job.setAlgorithm(m_pool.coin().algorithm(job.blob()[0]));
+    }
+
     m_job           = std::move(job);
     m_blocktemplate = std::move(blocktemplate);
     m_prevHash      = Json::getString(params, "prev_hash");