Merge df031be628 into 8afd4d5f2f

Fixed Zephyr mining (OpenCL)

Dockerfile

@@ -0,0 +1,45 @@
+# Build in disposable container so run-time container is small
+FROM alpine:latest as build
+
+# Build from master by default but allow build time specification
+ARG ref=master
+ENV my_ref=$ref
+
+# Developers may wish to specify an alternate repository for source
+ARG repo=https://github.com/xmrig/xmrig.git
+ENV my_repo=$repo
+
+RUN set -ex && \
+        # testing required for hwloc
+	echo @testing http://nl.alpinelinux.org/alpine/edge/testing >> /etc/apk/repositories
+
+RUN set -ex && \
+	apk --no-cache --update add \
+	    coreutils file grep openssl tar binutils \
+	    cmake g++ git linux-headers libpthread-stubs make hwloc-dev@testing \
+	    libuv-dev openssl-dev
+
+WORKDIR /usr/local/src
+
+RUN set -ex && \
+    	git clone $my_repo xmrig && \
+	cd xmrig && git checkout $my_ref && \
+	cmake -B build && \
+	cd build && \
+	make 
+
+# runtime container
+FROM alpine:latest
+
+RUN set -ex && \
+        # testing required for hwloc
+	echo @testing http://nl.alpinelinux.org/alpine/edge/testing >> /etc/apk/repositories
+
+RUN set -ex && \
+	apk --no-cache --update add \
+	# required libraries packages
+		openssl libuv hwloc@testing
+
+COPY --from=build /usr/local/src/xmrig/build/xmrig /bin/
+
+ENTRYPOINT ["/bin/xmrig"]

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

@@ -138,6 +138,93 @@ __kernel void blake2b_initial_hash(__global void *out, __global const void* bloc
 	t[7] = hash[7];
 }
 
+void blake2b_512_process_double_block_variable(ulong *out, ulong* m, __global const ulong* in, uint in_len, uint out_len)
+{
+	ulong v[16] =
+	{
+		iv0 ^ (0x01010000u | out_len), iv1, iv2, iv3, iv4      , iv5, iv6, iv7,
+		iv0               , iv1, iv2, iv3, iv4 ^ 128, iv5, iv6, iv7,
+	};
+
+	BLAKE2B_ROUNDS();
+
+	ulong h[8];
+	v[0] = h[0] = v[0] ^ v[8] ^ iv0 ^ (0x01010000u | out_len);
+	v[1] = h[1] = v[1] ^ v[9] ^ iv1;
+	v[2] = h[2] = v[2] ^ v[10] ^ iv2;
+	v[3] = h[3] = v[3] ^ v[11] ^ iv3;
+	v[4] = h[4] = v[4] ^ v[12] ^ iv4;
+	v[5] = h[5] = v[5] ^ v[13] ^ iv5;
+	v[6] = h[6] = v[6] ^ v[14] ^ iv6;
+	v[7] = h[7] = v[7] ^ v[15] ^ iv7;
+	v[8] = iv0;
+	v[9] = iv1;
+	v[10] = iv2;
+	v[11] = iv3;
+	v[12] = iv4 ^ in_len;
+	v[13] = iv5;
+	v[14] = ~iv6;
+	v[15] = iv7;
+
+	m[ 0] = (in_len > 128) ? in[16] : 0;
+	m[ 1] = (in_len > 136) ? in[17] : 0;
+	m[ 2] = (in_len > 144) ? in[18] : 0;
+	m[ 3] = (in_len > 152) ? in[19] : 0;
+	m[ 4] = (in_len > 160) ? in[20] : 0;
+	m[ 5] = (in_len > 168) ? in[21] : 0;
+	m[ 6] = (in_len > 176) ? in[22] : 0;
+	m[ 7] = (in_len > 184) ? in[23] : 0;
+	m[ 8] = (in_len > 192) ? in[24] : 0;
+	m[ 9] = (in_len > 200) ? in[25] : 0;
+	m[10] = (in_len > 208) ? in[26] : 0;
+	m[11] = (in_len > 216) ? in[27] : 0;
+	m[12] = (in_len > 224) ? in[28] : 0;
+	m[13] = (in_len > 232) ? in[29] : 0;
+	m[14] = (in_len > 240) ? in[30] : 0;
+	m[15] = (in_len > 248) ? in[31] : 0;
+
+	if (in_len % sizeof(ulong))
+		m[(in_len - 128) / sizeof(ulong)] &= (ulong)(-1) >> (64 - (in_len % sizeof(ulong)) * 8);
+
+	BLAKE2B_ROUNDS();
+
+	if (out_len >  0) out[0] = h[0] ^ v[0] ^ v[8];
+	if (out_len >  8) out[1] = h[1] ^ v[1] ^ v[9];
+	if (out_len > 16) out[2] = h[2] ^ v[2] ^ v[10];
+	if (out_len > 24) out[3] = h[3] ^ v[3] ^ v[11];
+	if (out_len > 32) out[4] = h[4] ^ v[4] ^ v[12];
+	if (out_len > 40) out[5] = h[5] ^ v[5] ^ v[13];
+	if (out_len > 48) out[6] = h[6] ^ v[6] ^ v[14];
+	if (out_len > 56) out[7] = h[7] ^ v[7] ^ v[15];
+}
+
+__attribute__((reqd_work_group_size(64, 1, 1)))
+__kernel void blake2b_initial_hash_double(__global void *out, __global const void* blockTemplate, uint blockTemplateSize, uint start_nonce)
+{
+	const uint global_index = get_global_id(0);
+
+	__global const ulong* p = (__global const ulong*) blockTemplate;
+
+	ulong m[16] = { p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], p[8], p[9], p[10], p[11], p[12], p[13], p[14], p[15] };
+
+	const ulong nonce = start_nonce + global_index;
+	m[4] = (m[4] & ((ulong)(-1) >>  8)) | (nonce << 56);
+	m[5] = (m[5] & ((ulong)(-1) << 24)) | (nonce >>  8);
+
+	ulong hash[8];
+	blake2b_512_process_double_block_variable(hash, m, p, blockTemplateSize, 64);
+
+	__global ulong* t = ((__global ulong*) out) + global_index * 8;
+	t[0] = hash[0];
+	t[1] = hash[1];
+	t[2] = hash[2];
+	t[3] = hash[3];
+	t[4] = hash[4];
+	t[5] = hash[5];
+	t[6] = hash[6];
+	t[7] = hash[7];
+}
+
 #define in_len 256
 
 #define out_len 32

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

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

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

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

src/backend/opencl/opencl.cmake

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

src/backend/opencl/runners/OclRxBaseRunner.cpp

@@ -25,6 +25,7 @@
 #include "backend/opencl/runners/OclRxBaseRunner.h"
 #include "backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h"
 #include "backend/opencl/kernels/rx/Blake2bInitialHashKernel.h"
+#include "backend/opencl/kernels/rx/Blake2bInitialHashDoubleKernel.h"
 #include "backend/opencl/kernels/rx/FillAesKernel.h"
 #include "backend/opencl/kernels/rx/FindSharesKernel.h"
 #include "backend/opencl/kernels/rx/HashAesKernel.h"
@@ -71,6 +72,7 @@ xmrig::OclRxBaseRunner::~OclRxBaseRunner()
     delete m_fillAes4Rx4_entropy;
     delete m_hashAes1Rx4;
     delete m_blake2b_initial_hash;
+    delete m_blake2b_initial_hash_double;
     delete m_blake2b_hash_registers_32;
     delete m_blake2b_hash_registers_64;
     delete m_find_shares;
@@ -87,12 +89,28 @@ void xmrig::OclRxBaseRunner::run(uint32_t nonce, uint32_t *hashOutput)
 {
     static const uint32_t zero = 0;
 
-    m_blake2b_initial_hash->setNonce(nonce);
+    if (m_jobSize <= 128) {
+        m_blake2b_initial_hash->setNonce(nonce);
+    }
+    else if (m_jobSize <= 256) {
+        m_blake2b_initial_hash_double->setNonce(nonce);
+    }
+    else {
+        hashOutput[0xFF] = 0;
+        return;
+    }
+
     m_find_shares->setNonce(nonce);
 
     enqueueWriteBuffer(m_output, CL_FALSE, sizeof(cl_uint) * 0xFF, sizeof(uint32_t), &zero);
 
-    m_blake2b_initial_hash->enqueue(m_queue, m_intensity);
+    if (m_jobSize <= 128) {
+        m_blake2b_initial_hash->enqueue(m_queue, m_intensity);
+    }
+    else {
+        m_blake2b_initial_hash_double->enqueue(m_queue, m_intensity);
+    }
+
     m_fillAes1Rx4_scratchpad->enqueue(m_queue, m_intensity);
 
     const uint32_t programCount = RxAlgo::programCount(m_algorithm);
@@ -134,7 +152,11 @@ void xmrig::OclRxBaseRunner::set(const Job &job, uint8_t *blob)
 
     enqueueWriteBuffer(m_input, CL_TRUE, 0, Job::kMaxBlobSize, blob);
 
+    m_jobSize = job.size();
+
     m_blake2b_initial_hash->setBlobSize(job.size());
+    m_blake2b_initial_hash_double->setBlobSize(job.size());
+
     m_find_shares->setTarget(job.target());
 }
 
@@ -166,6 +188,9 @@ void xmrig::OclRxBaseRunner::build()
     m_blake2b_initial_hash = new Blake2bInitialHashKernel(m_program);
     m_blake2b_initial_hash->setArgs(m_hashes, m_input);
 
+    m_blake2b_initial_hash_double = new Blake2bInitialHashDoubleKernel(m_program);
+    m_blake2b_initial_hash_double->setArgs(m_hashes, m_input);
+
     m_blake2b_hash_registers_32 = new Blake2bHashRegistersKernel(m_program, "blake2b_hash_registers_32");
     m_blake2b_hash_registers_64 = new Blake2bHashRegistersKernel(m_program, "blake2b_hash_registers_64");
 

src/backend/opencl/runners/OclRxBaseRunner.h

@@ -35,6 +35,7 @@ namespace xmrig {
 
 class Blake2bHashRegistersKernel;
 class Blake2bInitialHashKernel;
+class Blake2bInitialHashDoubleKernel;
 class FillAesKernel;
 class FindSharesKernel;
 class HashAesKernel;
@@ -58,21 +59,24 @@ protected:
 protected:
     virtual void execute(uint32_t iteration) = 0;
 
-    Blake2bHashRegistersKernel *m_blake2b_hash_registers_32 = nullptr;
-    Blake2bHashRegistersKernel *m_blake2b_hash_registers_64 = nullptr;
-    Blake2bInitialHashKernel *m_blake2b_initial_hash        = nullptr;
+    Blake2bHashRegistersKernel *m_blake2b_hash_registers_32       = nullptr;
+    Blake2bHashRegistersKernel *m_blake2b_hash_registers_64       = nullptr;
+    Blake2bInitialHashKernel *m_blake2b_initial_hash              = nullptr;
+    Blake2bInitialHashDoubleKernel *m_blake2b_initial_hash_double = nullptr;
     Buffer m_seed;
-    cl_mem m_dataset                                        = nullptr;
-    cl_mem m_entropy                                        = nullptr;
-    cl_mem m_hashes                                         = nullptr;
-    cl_mem m_rounding                                       = nullptr;
-    cl_mem m_scratchpads                                    = nullptr;
-    FillAesKernel *m_fillAes1Rx4_scratchpad                 = nullptr;
-    FillAesKernel *m_fillAes4Rx4_entropy                    = nullptr;
-    FindSharesKernel *m_find_shares                         = nullptr;
-    HashAesKernel *m_hashAes1Rx4                            = nullptr;
-    uint32_t m_gcn_version                                  = 12;
-    uint32_t m_worksize                                     = 8;
+    cl_mem m_dataset                                              = nullptr;
+    cl_mem m_entropy                                              = nullptr;
+    cl_mem m_hashes                                               = nullptr;
+    cl_mem m_rounding                                             = nullptr;
+    cl_mem m_scratchpads                                          = nullptr;
+    FillAesKernel *m_fillAes1Rx4_scratchpad                       = nullptr;
+    FillAesKernel *m_fillAes4Rx4_entropy                          = nullptr;
+    FindSharesKernel *m_find_shares                               = nullptr;
+    HashAesKernel *m_hashAes1Rx4                                  = nullptr;
+    uint32_t m_gcn_version                                        = 12;
+    uint32_t m_worksize                                           = 8;
+
+    size_t m_jobSize                                              = 0;
 };
 
 

src/base/net/stratum/Pool.h

@@ -1,7 +1,7 @@
 /* XMRig
  * Copyright (c) 2019      Howard Chu  <https://github.com/hyc>
- * Copyright (c) 2018-2021 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2021 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright (c) 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright (c) 2016-2024 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
  *
  *   This program is free software: you can redistribute it and/or modify
  *   it under the terms of the GNU General Public License as published by
@@ -59,7 +59,7 @@ public:
     static const char *kCoin;
     static const char *kDaemon;
     static const char *kDaemonPollInterval;
-    static const char* kDaemonJobTimeout;
+    static const char *kDaemonJobTimeout;
     static const char *kEnabled;
     static const char *kFingerprint;
     static const char *kKeepalive;
@@ -70,11 +70,11 @@ public:
     static const char *kSOCKS5;
     static const char *kSubmitToOrigin;
     static const char *kTls;
-    static const char* kSni;
+    static const char *kSni;
     static const char *kUrl;
     static const char *kUser;
-    static const char* kSpendSecretKey;
-    static const char* kDaemonZMQPort;
+    static const char *kSpendSecretKey;
+    static const char *kDaemonZMQPort;
     static const char *kNicehashHost;
 
     constexpr static int kKeepAliveTimeout         = 60;

src/core/config/usage.h

@@ -4,8 +4,8 @@
  * Copyright (c) 2014      Lucas Jones <https://github.com/lucasjones>
  * Copyright (c) 2014-2016 Wolf9466    <https://github.com/OhGodAPet>
  * Copyright (c) 2016      Jay D Dee   <jayddee246@gmail.com>
- * Copyright (c) 2018-2021 SChernykh   <https://github.com/SChernykh>
- * Copyright (c) 2016-2021 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ * Copyright (c) 2018-2024 SChernykh   <https://github.com/SChernykh>
+ * Copyright (c) 2016-2024 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
  *
  *   This program is free software: you can redistribute it and/or modify
  *   it under the terms of the GNU General Public License as published by
@@ -64,7 +64,7 @@ static inline const std::string &usage()
 
 #   ifdef XMRIG_FEATURE_HTTP
     u += "      --daemon                  use daemon RPC instead of pool for solo mining\n";
-    u += "      --daemon-zmq-port         daemon's zmq-pub port number (only use it if daemon has it enabled)\n";
+    u += "      --daemon-zmq-port=N       daemon's zmq-pub port number (only use it if daemon has it enabled)\n";
     u += "      --daemon-poll-interval=N  daemon poll interval in milliseconds (default: 1000)\n";
     u += "      --daemon-job-timeout=N    daemon job timeout in milliseconds (default: 15000)\n";
     u += "      --self-select=URL         self-select block templates from URL\n";