xmrig · Mar 15, 2022 · Mar 14, 2022
diff --git a/scripts/generate_cl.js b/scripts/generate_cl.js
@@ -76,6 +76,14 @@ function astrobwt()
 }
 
 
+function astrobwt_v2()
+{
+    const astrobwt = opencl_minify(addIncludes('astrobwt_v2.cl', [ 'BWT.cl', 'salsa20.cl', 'sha3.cl' ]));
+
+    fs.writeFileSync('astrobwt_v2_cl.h', text2h(astrobwt, 'xmrig', 'astrobwt_v2_cl'));
+}
+
+
 function kawpow()
 {
     const kawpow = opencl_minify(addIncludes('kawpow.cl', [ 'defs.h' ]));
@@ -102,6 +110,11 @@ process.chdir(path.resolve('src/backend/opencl/cl/astrobwt'));
 
 astrobwt();
 
+process.chdir(cwd);
+process.chdir(path.resolve('src/backend/opencl/cl/astrobwt_v2'));
+
+astrobwt_v2();
+
 process.chdir(cwd);
 process.chdir(path.resolve('src/backend/opencl/cl/kawpow'));
 

diff --git a/src/backend/cuda/CudaBackend.cpp b/src/backend/cuda/CudaBackend.cpp
@@ -221,7 +221,7 @@ class CudaBackendPrivate
         size_t algo_l3 = algo.l3();
 
 #       ifdef XMRIG_ALGO_ASTROBWT
-        if (algo.family() == Algorithm::ASTROBWT) {
+        if (algo.id() == Algorithm::ASTROBWT_DERO) {
             algo_l3 = CudaAstroBWTRunner::BWT_DATA_STRIDE * 17 + 1024;
         }
 #       endif

diff --git a/src/backend/opencl/OclBackend.cpp b/src/backend/opencl/OclBackend.cpp
@@ -205,7 +205,7 @@ class OclBackendPrivate
         size_t algo_l3 = algo.l3();
 
 #       ifdef XMRIG_ALGO_ASTROBWT
-        if (algo.family() == Algorithm::ASTROBWT) {
+        if (algo.id() == Algorithm::ASTROBWT_DERO) {
             algo_l3 = OclAstroBWTRunner::BWT_DATA_STRIDE * 17 + 324;
         }
 #       endif

diff --git a/src/backend/opencl/OclConfig_gen.h b/src/backend/opencl/OclConfig_gen.h
@@ -133,6 +133,12 @@ size_t inline generate<Algorithm::RANDOM_X>(Threads<OclThreads> &threads, const
 template<>
 size_t inline generate<Algorithm::ASTROBWT>(Threads<OclThreads>& threads, const std::vector<OclDevice>& devices)
 {
+    size_t count = 0;
+
+    if (!threads.isExist(Algorithm::ASTROBWT_DERO_2)) {
+        count += threads.move(Algorithm::kASTROBWT_DERO_2, OclThreads(devices, Algorithm::ASTROBWT_DERO_2));
+    }
+
     return generate(Algorithm::kASTROBWT, threads, Algorithm::ASTROBWT_DERO, devices);
 }
 #endif

diff --git a/src/backend/opencl/OclWorker.cpp b/src/backend/opencl/OclWorker.cpp
@@ -37,6 +37,7 @@
 
 #ifdef XMRIG_ALGO_ASTROBWT
 #   include "backend/opencl/runners/OclAstroBWTRunner.h"
+#   include "backend/opencl/runners/OclAstroBWT_v2_Runner.h"
 #endif
 
 #ifdef XMRIG_ALGO_KAWPOW
@@ -92,7 +93,12 @@ xmrig::OclWorker::OclWorker(size_t id, const OclLaunchData &data) :
 
     case Algorithm::ASTROBWT:
 #       ifdef XMRIG_ALGO_ASTROBWT
-        m_runner = new OclAstroBWTRunner(id, data);
+        if (m_algorithm.id() == Algorithm::ASTROBWT_DERO_2) {
+            m_runner = new OclAstroBWT_v2_Runner(id, data);
+        }
+        else {
+            m_runner = new OclAstroBWTRunner(id, data);
+        }
 #       endif
         break;
 

diff --git a/src/backend/opencl/cl/OclSource.cpp b/src/backend/opencl/cl/OclSource.cpp
@@ -34,6 +34,7 @@
 
 #ifdef XMRIG_ALGO_ASTROBWT
 #   include "backend/opencl/cl/astrobwt/astrobwt_cl.h"
+#   include "backend/opencl/cl/astrobwt_v2/astrobwt_v2_cl.h"
 #endif
 
 #ifdef XMRIG_ALGO_KAWPOW
@@ -52,7 +53,7 @@ const char *xmrig::OclSource::get(const Algorithm &algorithm)
 
 #   ifdef XMRIG_ALGO_ASTROBWT
     if (algorithm.family() == Algorithm::ASTROBWT) {
-        return astrobwt_cl;
+        return (algorithm.id() == Algorithm::ASTROBWT_DERO_2) ? astrobwt_v2_cl : astrobwt_cl;
     }
 #   endif
 

diff --git a/src/backend/opencl/cl/astrobwt_v2/BWT.cl b/src/backend/opencl/cl/astrobwt_v2/BWT.cl
@@ -0,0 +1,179 @@
+/* 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-2022 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2022 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/>.
+ */
+
+#define BLOCK_SIZE 1024
+#define DATA_SIZE 9973
+#define DATA_STRIDE 10240
+#define BITS 14
+#define COUNTERS_SIZE (1 << BITS)
+
+inline uint16_t atomic_inc16(__local uint16_t* value)
+{
+	const size_t k = (size_t) value;
+	if ((k & 3) == 0) {
+		return atomic_add((__local uint32_t*) value, 1);
+	}
+	return atomic_add((__local uint32_t*)(k - 2), 0x10000) >> 16;
+}
+
+__attribute__((reqd_work_group_size(BLOCK_SIZE, 1, 1)))
+__kernel void BWT_preprocess(__global const uint8_t* datas, __global uint32_t* keys)
+{
+	const uint32_t data_offset = get_group_id(0) * DATA_STRIDE;
+	const uint32_t tid = get_local_id(0);
+
+	__local uint32_t counters_buf[COUNTERS_SIZE / 2];
+	__local uint16_t* counters = (__local uint16_t*) counters_buf;
+	for (uint32_t i = tid; i < COUNTERS_SIZE / 2; i += BLOCK_SIZE) {
+		counters_buf[i] = 0;
+	}
+	barrier(CLK_LOCAL_MEM_FENCE);
+
+	datas += data_offset;
+	keys += data_offset;
+
+	for (uint32_t i = tid; i < DATA_SIZE; i += BLOCK_SIZE) {
+		const uint32_t k0 = datas[i];
+		const uint32_t k1 = datas[i + 1];
+		const uint32_t k = ((k0 << 8) | k1) >> (16 - BITS);
+		atomic_inc16(counters + k);
+	}
+	barrier(CLK_LOCAL_MEM_FENCE);
+
+	#pragma unroll BITS
+	for (int k = 0; k < BITS; ++k) {
+		for (uint32_t t1 = tid; t1 < ((COUNTERS_SIZE / 2) >> k); t1 += BLOCK_SIZE) {
+			const uint32_t i = (t1 << (k + 1)) + ((1 << (k + 1)) - 1);
+			counters[i] += counters[i - (1 << k)];
+		}
+		barrier(CLK_LOCAL_MEM_FENCE);
+	}
+
+	if (tid == 0) {
+		counters[COUNTERS_SIZE - 1] = 0;
+	}
+	barrier(CLK_LOCAL_MEM_FENCE);
+
+	#pragma unroll BITS
+	for (int k = BITS - 1; k >= 0; --k) {
+		for (uint32_t t1 = tid; t1 < ((COUNTERS_SIZE / 2) >> k); t1 += BLOCK_SIZE) {
+			const uint32_t i = (t1 << (k + 1)) + ((1 << (k + 1)) - 1);
+			const uint16_t old = counters[i];
+			counters[i] = old + counters[i - (1 << k)];
+			counters[i - (1 << k)] = old;
+		}
+		barrier(CLK_LOCAL_MEM_FENCE);
+	}
+
+	for (uint32_t i = tid; i < DATA_SIZE; i += BLOCK_SIZE) {
+		const uint32_t k0 = datas[i];
+		const uint32_t k1 = datas[i + 1];
+		const uint32_t k = (k0 << 8) | k1;
+		const uint32_t index = atomic_inc16(counters + (k >> (16 - BITS)));
+		keys[index] = (k << 16) | i;
+	}
+}
+
+inline void fix_order(__global const uint8_t* input, uint32_t a, uint32_t b, __global uint32_t* keys)
+{
+	const uint32_t ka = keys[a];
+	const uint32_t kb = keys[b];
+	const uint32_t index_a = ka & 0xFFFF;
+	const uint32_t index_b = kb & 0xFFFF;
+
+	const uint32_t value_a =
+		(((uint32_t)input[index_a + 1]) << 24) |
+		(((uint32_t)input[index_a + 2]) << 16) |
+		(((uint32_t)input[index_a + 3]) << 8) |
+		((uint32_t)input[index_a + 4]);
+
+	const uint32_t value_b =
+		(((uint32_t)input[index_b + 1]) << 24) |
+		(((uint32_t)input[index_b + 2]) << 16) |
+		(((uint32_t)input[index_b + 3]) << 8) |
+		((uint32_t)input[index_b + 4]);
+
+	if (value_a > value_b)
+	{
+		keys[a] = kb;
+		keys[b] = ka;
+	}
+}
+
+__attribute__((reqd_work_group_size(BLOCK_SIZE, 1, 1)))
+__kernel void BWT_fix_order(__global const uint8_t* datas, __global uint32_t* keys, __global uint16_t* values)
+{
+	const uint32_t tid = get_local_id(0);
+	const uint32_t gid = get_group_id(0);
+
+	const uint32_t data_offset = gid * 10240;
+
+	const uint32_t N = 9973;
+
+	datas += data_offset;
+	keys += data_offset;
+	values += data_offset;
+
+	for (uint32_t i = tid, N1 = N - 1; i < N1; i += BLOCK_SIZE)
+	{
+		const uint32_t value = keys[i] >> (32 - BITS);
+		if (value == (keys[i + 1] >> (32 - BITS)))
+		{
+			if (i && (value == (keys[i - 1] >> (32 - BITS))))
+				continue;
+
+			uint32_t n = i + 2;
+			while ((n < N) && (value == (keys[n] >> (32 - BITS))))
+				++n;
+
+			for (uint32_t j = i; j < n; ++j)
+				for (uint32_t k = j + 1; k < n; ++k)
+					fix_order(datas, j, k, keys);
+		}
+	}
+	barrier(CLK_GLOBAL_MEM_FENCE);
+
+	for (uint32_t i = tid; i < N; i += BLOCK_SIZE) {
+		values[i] = keys[i];
+	}
+}
+
+__kernel void find_shares(__global const uint64_t* hashes, uint64_t target, __global uint32_t* shares)
+{
+	const uint32_t global_index = get_global_id(0);
+
+	if (hashes[global_index * 4 + 3] >= target) {
+		return;
+	}
+
+	const uint32_t idx = atomic_inc(shares + 0xFF);
+	if (idx < 0xFF)
+		shares[idx] = global_index;
+}
+
+#undef BLOCK_SIZE
+#undef DATA_SIZE
+#undef DATA_STRIDE
+#undef BITS
+#undef COUNTERS_SIZE
diff --git a/src/backend/opencl/cl/astrobwt_v2/astrobwt_v2.cl b/src/backend/opencl/cl/astrobwt_v2/astrobwt_v2.cl
@@ -0,0 +1,35 @@
+/* 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-2022 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2022 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/>.
+ */
+
+typedef uchar uint8_t;
+typedef ushort uint16_t;
+typedef uint uint32_t;
+typedef ulong uint64_t;
+
+typedef int int32_t;
+typedef long int64_t;
+
+#include "BWT.cl"
+#include "salsa20.cl"
+#include "sha3.cl"