From fbedf197aba11b24b3389818ffa45041d34d72d7 Mon Sep 17 00:00:00 2001
From: SChernykh <sergey.v.chernykh@gmail.com>
Date: Sun, 22 Mar 2020 22:36:21 +0100
Subject: [PATCH] AstroBWT OpenCL code

---
 scripts/generate_cl.js                        |  16 +-
 src/backend/opencl/OclBackend.cpp             |  11 +-
 src/backend/opencl/OclConfig.cpp              |   1 +
 src/backend/opencl/OclConfig_gen.h            |   9 +
 src/backend/opencl/OclThread.cpp              |   6 +-
 src/backend/opencl/OclThread.h                |  15 +
 src/backend/opencl/OclWorker.cpp              |  16 +-
 src/backend/opencl/cl/OclSource.cpp           |  10 +
 src/backend/opencl/cl/astrobwt/BWT.cl         | 199 +++++++++
 src/backend/opencl/cl/astrobwt/astrobwt.cl    |  35 ++
 src/backend/opencl/cl/astrobwt/astrobwt_cl.h  | 395 ++++++++++++++++++
 src/backend/opencl/cl/astrobwt/salsa20.cl     | 153 +++++++
 src/backend/opencl/cl/astrobwt/sha3.cl        | 198 +++++++++
 .../ocl_generic_astrobwt_generator.cpp        |  64 +++
 src/backend/opencl/interfaces/IOclRunner.h    |   2 +
 .../astrobwt/AstroBWT_FilterKernel.cpp        |  42 ++
 .../kernels/astrobwt/AstroBWT_FilterKernel.h  |  48 +++
 .../astrobwt/AstroBWT_FindSharesKernel.cpp    |  47 +++
 .../astrobwt/AstroBWT_FindSharesKernel.h      |  49 +++
 .../kernels/astrobwt/AstroBWT_MainKernel.cpp  |  44 ++
 .../kernels/astrobwt/AstroBWT_MainKernel.h    |  48 +++
 .../astrobwt/AstroBWT_PrepareBatch2Kernel.cpp |  41 ++
 .../astrobwt/AstroBWT_PrepareBatch2Kernel.h   |  48 +++
 .../astrobwt/AstroBWT_SHA3InitialKernel.cpp   |  44 ++
 .../astrobwt/AstroBWT_SHA3InitialKernel.h     |  48 +++
 .../kernels/astrobwt/AstroBWT_SHA3Kernel.cpp  |  44 ++
 .../kernels/astrobwt/AstroBWT_SHA3Kernel.h    |  48 +++
 .../astrobwt/AstroBWT_Salsa20Kernel.cpp       |  43 ++
 .../kernels/astrobwt/AstroBWT_Salsa20Kernel.h |  48 +++
 src/backend/opencl/opencl.cmake               |  25 ++
 .../opencl/runners/OclAstroBWTRunner.cpp      | 211 ++++++++++
 .../opencl/runners/OclAstroBWTRunner.h        |  95 +++++
 src/backend/opencl/runners/OclBaseRunner.h    |   2 +
 src/backend/opencl/wrappers/OclDevice.cpp     |   7 +
 34 files changed, 2106 insertions(+), 6 deletions(-)
 create mode 100644 src/backend/opencl/cl/astrobwt/BWT.cl
 create mode 100644 src/backend/opencl/cl/astrobwt/astrobwt.cl
 create mode 100644 src/backend/opencl/cl/astrobwt/astrobwt_cl.h
 create mode 100644 src/backend/opencl/cl/astrobwt/salsa20.cl
 create mode 100644 src/backend/opencl/cl/astrobwt/sha3.cl
 create mode 100644 src/backend/opencl/generators/ocl_generic_astrobwt_generator.cpp
 create mode 100644 src/backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.cpp
 create mode 100644 src/backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.h
 create mode 100644 src/backend/opencl/kernels/astrobwt/AstroBWT_FindSharesKernel.cpp
 create mode 100644 src/backend/opencl/kernels/astrobwt/AstroBWT_FindSharesKernel.h
 create mode 100644 src/backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.cpp
 create mode 100644 src/backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.h
 create mode 100644 src/backend/opencl/kernels/astrobwt/AstroBWT_PrepareBatch2Kernel.cpp
 create mode 100644 src/backend/opencl/kernels/astrobwt/AstroBWT_PrepareBatch2Kernel.h
 create mode 100644 src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3InitialKernel.cpp
 create mode 100644 src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3InitialKernel.h
 create mode 100644 src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.cpp
 create mode 100644 src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.h
 create mode 100644 src/backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.cpp
 create mode 100644 src/backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.h
 create mode 100644 src/backend/opencl/runners/OclAstroBWTRunner.cpp
 create mode 100644 src/backend/opencl/runners/OclAstroBWTRunner.h

diff --git a/scripts/generate_cl.js b/scripts/generate_cl.js
index 453a50fd3..fb2a387a6 100644
--- a/scripts/generate_cl.js
+++ b/scripts/generate_cl.js
@@ -76,6 +76,15 @@ function rx()
 }
 
 
+function astrobwt()
+{
+    const astrobwt = opencl_minify(addIncludes('astrobwt.cl', [ 'BWT.cl', 'salsa20.cl', 'sha3.cl' ]));
+
+    // fs.writeFileSync('astrobwt_gen.cl', astrobwt);
+    fs.writeFileSync('astrobwt_cl.h', text2h(astrobwt, 'xmrig', 'astrobwt_cl'));
+}
+
+
 process.chdir(path.resolve('src/backend/opencl/cl/cn'));
 
 cn();
@@ -85,4 +94,9 @@ cn_gpu();
 process.chdir(cwd);
 process.chdir(path.resolve('src/backend/opencl/cl/rx'));
 
-rx();
\ No newline at end of file
+rx();
+
+process.chdir(cwd);
+process.chdir(path.resolve('src/backend/opencl/cl/astrobwt'));
+
+astrobwt();
diff --git a/src/backend/opencl/OclBackend.cpp b/src/backend/opencl/OclBackend.cpp
index 5cbf57b70..26580e411 100644
--- a/src/backend/opencl/OclBackend.cpp
+++ b/src/backend/opencl/OclBackend.cpp
@@ -36,6 +36,7 @@
 #include "backend/opencl/OclLaunchData.h"
 #include "backend/opencl/OclWorker.h"
 #include "backend/opencl/runners/tools/OclSharedState.h"
+#include "backend/opencl/runners/OclAstroBWTRunner.h"
 #include "backend/opencl/wrappers/OclContext.h"
 #include "backend/opencl/wrappers/OclLib.h"
 #include "base/io/log/Log.h"
@@ -202,6 +203,14 @@ public:
 
         Log::print(WHITE_BOLD("|  # | GPU |  BUS ID |    I |  W | SI | MC |  U |  MEM | NAME"));
 
+        size_t algo_l3 = algo.l3();
+
+#       ifdef XMRIG_ALGO_ASTROBWT
+        if (algo.family() == Algorithm::ASTROBWT) {
+            algo_l3 = xmrig::OclAstroBWTRunner::BWT_DATA_STRIDE * 17 + 324;
+        }
+#       endif
+
         size_t i = 0;
         for (const auto &data : threads) {
             Log::print("|" CYAN_BOLD("%3zu") " |" CYAN_BOLD("%4u") " |" YELLOW(" %7s") " |" CYAN_BOLD("%5u") " |" CYAN_BOLD("%3u") " |"
@@ -214,7 +223,7 @@ public:
                        data.thread.stridedIndex(),
                        data.thread.stridedIndex() == 2 ? std::to_string(data.thread.memChunk()).c_str() : "-",
                        data.thread.unrollFactor(),
-                       data.thread.intensity() * algo.l3() / oneMiB,
+                       data.thread.intensity() * algo_l3 / oneMiB,
                        data.device.printableName().data()
                        );
 
diff --git a/src/backend/opencl/OclConfig.cpp b/src/backend/opencl/OclConfig.cpp
index a90bcc903..b06746f3c 100644
--- a/src/backend/opencl/OclConfig.cpp
+++ b/src/backend/opencl/OclConfig.cpp
@@ -220,6 +220,7 @@ void xmrig::OclConfig::generate()
     count += xmrig::generate<Algorithm::CN_HEAVY>(m_threads, devices);
     count += xmrig::generate<Algorithm::CN_PICO>(m_threads, devices);
     count += xmrig::generate<Algorithm::RANDOM_X>(m_threads, devices);
+    count += xmrig::generate<Algorithm::ASTROBWT>(m_threads, devices);
 
     m_shouldSave = count > 0;
 }
diff --git a/src/backend/opencl/OclConfig_gen.h b/src/backend/opencl/OclConfig_gen.h
index 1c8a6a437..5feacc6fb 100644
--- a/src/backend/opencl/OclConfig_gen.h
+++ b/src/backend/opencl/OclConfig_gen.h
@@ -130,6 +130,15 @@ size_t inline generate<Algorithm::RANDOM_X>(Threads<OclThreads> &threads, const
 #endif
 
 
+#ifdef XMRIG_ALGO_ASTROBWT
+template<>
+size_t inline generate<Algorithm::ASTROBWT>(Threads<OclThreads>& threads, const std::vector<OclDevice>& devices)
+{
+    return generate("astrobwt", threads, Algorithm::ASTROBWT_DERO, devices);
+}
+#endif
+
+
 static inline std::vector<OclDevice> filterDevices(const std::vector<OclDevice> &devices, const std::vector<uint32_t> &hints)
 {
     std::vector<OclDevice> out;
diff --git a/src/backend/opencl/OclThread.cpp b/src/backend/opencl/OclThread.cpp
index e94eb876b..b542b443f 100644
--- a/src/backend/opencl/OclThread.cpp
+++ b/src/backend/opencl/OclThread.cpp
@@ -123,7 +123,9 @@ rapidjson::Value xmrig::OclThread::toJSON(rapidjson::Document &doc) const
 
     out.AddMember(StringRef(kIndex),        index(), allocator);
     out.AddMember(StringRef(kIntensity),    intensity(), allocator);
-    out.AddMember(StringRef(kWorksize),     worksize(), allocator);
+    if (!m_fields.test(ASTROBWT_FIELDS)) {
+        out.AddMember(StringRef(kWorksize), worksize(), allocator);
+    }
 
     if (m_fields.test(STRIDED_INDEX_FIELD)) {
         Value si(kArrayType);
@@ -149,7 +151,7 @@ rapidjson::Value xmrig::OclThread::toJSON(rapidjson::Document &doc) const
         out.AddMember(StringRef(kDatasetHost),  isDatasetHost(), allocator);
 #       endif
     }
-    else {
+    else if (!m_fields.test(ASTROBWT_FIELDS)) {
         out.AddMember(StringRef(kUnroll), unrollFactor(), allocator);
     }
 
diff --git a/src/backend/opencl/OclThread.h b/src/backend/opencl/OclThread.h
index 4febb7a03..dac57a4f1 100644
--- a/src/backend/opencl/OclThread.h
+++ b/src/backend/opencl/OclThread.h
@@ -81,6 +81,20 @@ public:
     }
 #   endif
 
+#   ifdef XMRIG_ALGO_ASTROBWT
+    OclThread(uint32_t index, uint32_t intensity, uint32_t threads) :
+        m_fields(4),
+        m_threads(threads, -1),
+        m_index(index),
+        m_memChunk(0),
+        m_stridedIndex(0),
+        m_unrollFactor(1),
+        m_worksize(1)
+    {
+        setIntensity(intensity);
+    }
+#   endif
+
     OclThread(const rapidjson::Value &value);
 
     inline bool isAsm() const                               { return m_gcnAsm; }
@@ -105,6 +119,7 @@ private:
     enum Fields {
         STRIDED_INDEX_FIELD,
         RANDOMX_FIELDS,
+        ASTROBWT_FIELDS,
         FIELD_MAX
     };
 
diff --git a/src/backend/opencl/OclWorker.cpp b/src/backend/opencl/OclWorker.cpp
index c8f69b27f..10e0ab338 100644
--- a/src/backend/opencl/OclWorker.cpp
+++ b/src/backend/opencl/OclWorker.cpp
@@ -41,6 +41,10 @@
 #   include "backend/opencl/runners/OclRxVmRunner.h"
 #endif
 
+#ifdef XMRIG_ALGO_ASTROBWT
+#   include "backend/opencl/runners/OclAstroBWTRunner.h"
+#endif
+
 #ifdef XMRIG_ALGO_CN_GPU
 #   include "backend/opencl/runners/OclRyoRunner.h"
 #endif
@@ -96,6 +100,12 @@ xmrig::OclWorker::OclWorker(size_t id, const OclLaunchData &data) :
 #       endif
         break;
 
+    case Algorithm::ASTROBWT:
+#       ifdef XMRIG_ALGO_ASTROBWT
+        m_runner = new OclAstroBWTRunner(id, data);
+#       endif
+        break;
+
     default:
 #       ifdef XMRIG_ALGO_CN_GPU
         if (m_algorithm == Algorithm::CN_GPU) {
@@ -148,6 +158,8 @@ void xmrig::OclWorker::start()
 {
     cl_uint results[0x100];
 
+    const uint32_t runnerRoundSize = m_runner->roundSize();
+
     while (Nonce::sequence(Nonce::OPENCL) > 0) {
         if (!isReady()) {
             m_sharedData.setResumeCounter(0);
@@ -186,7 +198,7 @@ void xmrig::OclWorker::start()
                 JobResults::submit(m_job.currentJob(), results, results[0xFF]);
             }
 
-            if (!m_job.nextRound(roundSize(m_intensity), m_intensity)) {
+            if (!m_job.nextRound(roundSize(runnerRoundSize), runnerRoundSize)) {
                 JobResults::done(m_job.currentJob());
             }
 
@@ -228,7 +240,7 @@ void xmrig::OclWorker::storeStats(uint64_t t)
         return;
     }
 
-    m_count += m_intensity;
+    m_count += m_runner->processedHashes();
 
     m_sharedData.setRunTime(Chrono::steadyMSecs() - t);
 
diff --git a/src/backend/opencl/cl/OclSource.cpp b/src/backend/opencl/cl/OclSource.cpp
index b30e8a8f4..85ed4612a 100644
--- a/src/backend/opencl/cl/OclSource.cpp
+++ b/src/backend/opencl/cl/OclSource.cpp
@@ -36,6 +36,10 @@
 #   include "backend/opencl/cl/rx/randomx_cl.h"
 #endif
 
+#ifdef XMRIG_ALGO_ASTROBWT
+#   include "backend/opencl/cl/astrobwt/astrobwt_cl.h"
+#endif
+
 
 const char *xmrig::OclSource::get(const Algorithm &algorithm)
 {
@@ -45,6 +49,12 @@ const char *xmrig::OclSource::get(const Algorithm &algorithm)
     }
 #   endif
 
+#   ifdef XMRIG_ALGO_ASTROBWT
+    if (algorithm.family() == Algorithm::ASTROBWT) {
+        return astrobwt_cl;
+    }
+#   endif
+
 #   ifdef XMRIG_ALGO_CN_GPU
     if (algorithm == Algorithm::CN_GPU) {
         return cryptonight_gpu_cl;
diff --git a/src/backend/opencl/cl/astrobwt/BWT.cl b/src/backend/opencl/cl/astrobwt/BWT.cl
new file mode 100644
index 000000000..94551d954
--- /dev/null
+++ b/src/backend/opencl/cl/astrobwt/BWT.cl
@@ -0,0 +1,199 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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 STAGE1_SIZE 147253
+
+#define COUNTING_SORT_BITS 11
+#define COUNTING_SORT_SIZE (1 << COUNTING_SORT_BITS)
+#define FINAL_SORT_BATCH_SIZE COUNTING_SORT_SIZE
+#define FINAL_SORT_OVERLAP_SIZE 32
+
+__attribute__((reqd_work_group_size(BWT_GROUP_SIZE, 1, 1)))
+__kernel void BWT(__global uint8_t* datas, __global uint32_t* data_sizes, uint32_t data_stride, __global uint64_t* indices, __global uint64_t* tmp_indices)
+{
+	const uint32_t tid = get_local_id(0);
+	const uint32_t gid = get_group_id(0);
+
+	__local int counters[COUNTING_SORT_SIZE][2];
+
+	for (uint32_t i = tid; i < COUNTING_SORT_SIZE * 2; i += BWT_GROUP_SIZE)
+		((__local int*)counters)[i] = 0;
+
+	const uint64_t data_offset = (uint64_t)(gid) * data_stride;
+
+	__global uint8_t* input = datas + data_offset + 128;
+	const uint32_t N = data_sizes[gid] + 1;
+
+	__global uint64_t* p = (__global uint64_t*)(input);
+	volatile __local uint8_t* counters_atomic = (volatile __local uint8_t*)(counters);
+
+	indices += data_offset;
+	tmp_indices += data_offset;
+
+	for (uint32_t i = tid; i < N; i += BWT_GROUP_SIZE)
+	{
+		const uint32_t index = i >> 3;
+		const uint32_t bit_offset = (i & 7) << 3;
+
+		const uint64_t a = p[index];
+		uint64_t b = p[index + 1];
+		if (bit_offset == 0)
+			b = 0;
+
+		uint64_t value = (a >> bit_offset) | (b << (64 - bit_offset));
+
+		uint2 tmp;
+		const uchar4 mask = (uchar4)(3, 2, 1, 0);
+
+		tmp.x = as_uint(shuffle(as_uchar4(as_uint2(value).y), mask));
+		tmp.y = as_uint(shuffle(as_uchar4(as_uint2(value).x), mask));
+
+		value = as_ulong(tmp);
+
+		indices[i] = (value & ((uint64_t)(-1) << 21)) | i;
+		atomic_add((volatile __local int*)(counters_atomic + (((value >> (64 - COUNTING_SORT_BITS * 2)) & (COUNTING_SORT_SIZE - 1)) << 3)), 1);
+		atomic_add((volatile __local int*)(counters_atomic + ((value >> (64 - COUNTING_SORT_BITS)) << 3) + 4), 1);
+	}
+
+	if (tid == 0)
+	{
+		int t0 = counters[0][0];
+		int t1 = counters[0][1];
+		counters[0][0] = t0 - 1;
+		counters[0][1] = t1 - 1;
+		for (uint32_t i = 1; i < COUNTING_SORT_SIZE; ++i)
+		{
+			t0 += counters[i][0];
+			t1 += counters[i][1];
+			counters[i][0] = t0 - 1;
+			counters[i][1] = t1 - 1;
+		}
+	}
+
+	for (int i = tid; i < N; i += BWT_GROUP_SIZE)
+	{
+		const uint64_t data = indices[i];
+		const int k = atomic_sub((volatile __local int*)(counters_atomic + (((data >> (64 - COUNTING_SORT_BITS * 2)) & (COUNTING_SORT_SIZE - 1)) << 3)), 1);
+		tmp_indices[k] = data;
+	}
+
+	for (int i = N - 1 - tid; i >= 0; i -= BWT_GROUP_SIZE)
+	{
+		const uint64_t data = tmp_indices[i];
+		const int k = atomic_sub((volatile __local int*)(counters_atomic + ((data >> (64 - COUNTING_SORT_BITS)) << 3) + 4), 1);
+		indices[k] = data;
+	}
+
+	__local uint64_t* buf = (__local uint64_t*)(counters);
+	for (uint32_t i = 0; i < N; i += FINAL_SORT_BATCH_SIZE - FINAL_SORT_OVERLAP_SIZE)
+	{
+		const uint32_t len = (N - i < FINAL_SORT_BATCH_SIZE) ? (N - i) : FINAL_SORT_BATCH_SIZE;
+
+		for (uint32_t j = tid; j < len; j += BWT_GROUP_SIZE)
+			buf[j] = indices[i + j];
+
+		if (tid == 0)
+		{
+			uint64_t prev_t = buf[0];
+			for (int i = 1; i < len; ++i)
+			{
+				uint64_t t = buf[i];
+				if (t < prev_t)
+				{
+					const uint64_t t2 = prev_t;
+					int j = i - 1;
+					do
+					{
+						buf[j + 1] = prev_t;
+						--j;
+						if (j < 0)
+							break;
+						prev_t = buf[j];
+					} while (t < prev_t);
+					buf[j + 1] = t;
+					t = t2;
+				}
+				prev_t = t;
+			}
+		}
+
+		for (uint32_t j = tid; j < len; j += BWT_GROUP_SIZE)
+			indices[i + j] = buf[j];
+	}
+
+	--input;
+	__global uint8_t* output = (__global uint8_t*)(tmp_indices);
+	for (int i = tid; i <= N; i += BWT_GROUP_SIZE)
+		output[i] = input[indices[i] & ((1 << 21) - 1)];
+}
+
+__kernel void filter(uint32_t nonce, uint32_t bwt_max_size, __global const uint32_t* hashes, __global uint32_t* filtered_hashes)
+{
+	const uint32_t global_id = get_global_id(0);
+
+	__global const uint32_t* hash = hashes + global_id * (32 / sizeof(uint32_t));
+	const uint32_t stage2_size = STAGE1_SIZE + (*hash & 0xfffff);
+
+	if (stage2_size < bwt_max_size)
+	{
+		const int index = atomic_add((volatile __global int*)(filtered_hashes), 1) * (36 / sizeof(uint32_t)) + 1;
+
+		filtered_hashes[index] = nonce + global_id;
+
+		#pragma unroll(8)
+		for (uint32_t i = 0; i < 8; ++i)
+			filtered_hashes[index + i + 1] = hash[i];
+	}
+}
+
+__kernel void prepare_batch2(__global uint32_t* hashes, __global uint32_t* filtered_hashes, __global uint32_t* data_sizes)
+{
+	const uint32_t global_id = get_global_id(0);
+	const uint32_t N = filtered_hashes[0] - get_global_size(0);
+
+	if (global_id == 0)
+		filtered_hashes[0] = N;
+
+	__global uint32_t* hash = hashes + global_id * 8;
+	__global uint32_t* filtered_hash = filtered_hashes + (global_id + N) * 9 + 1;
+
+	const uint32_t stage2_size = STAGE1_SIZE + (filtered_hash[1] & 0xfffff);
+	data_sizes[global_id] = stage2_size;
+
+	#pragma unroll(8)
+	for (uint32_t i = 0; i < 8; ++i)
+		hash[i] = filtered_hash[i + 1];
+}
+
+__kernel void find_shares(__global const uint64_t* hashes, __global const uint32_t* filtered_hashes, uint64_t target, __global uint32_t* shares)
+{
+	const uint32_t global_index = get_global_id(0);
+
+	if (hashes[global_index * 4 + 3] < target)
+	{
+		const uint32_t idx = atomic_inc(shares + 0xFF);
+		if (idx < 0xFF)
+			shares[idx] = filtered_hashes[(filtered_hashes[0] + global_index) * 9 + 1];
+	}
+}
diff --git a/src/backend/opencl/cl/astrobwt/astrobwt.cl b/src/backend/opencl/cl/astrobwt/astrobwt.cl
new file mode 100644
index 000000000..476366c84
--- /dev/null
+++ b/src/backend/opencl/cl/astrobwt/astrobwt.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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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"
diff --git a/src/backend/opencl/cl/astrobwt/astrobwt_cl.h b/src/backend/opencl/cl/astrobwt/astrobwt_cl.h
new file mode 100644
index 000000000..200ac2449
--- /dev/null
+++ b/src/backend/opencl/cl/astrobwt/astrobwt_cl.h
@@ -0,0 +1,395 @@
+#pragma once
+
+namespace xmrig {
+
+static char astrobwt_cl[12378] = {
+    0x74,0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x75,0x63,0x68,0x61,0x72,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x3b,0x0a,0x74,0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x75,
+    0x73,0x68,0x6f,0x72,0x74,0x20,0x75,0x69,0x6e,0x74,0x31,0x36,0x5f,0x74,0x3b,0x0a,0x74,0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x75,0x69,0x6e,0x74,0x20,0x75,0x69,0x6e,
+    0x74,0x33,0x32,0x5f,0x74,0x3b,0x0a,0x74,0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x75,0x6c,0x6f,0x6e,0x67,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x3b,0x0a,0x74,
+    0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x69,0x6e,0x74,0x20,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x3b,0x0a,0x74,0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x6c,0x6f,0x6e,0x67,
+    0x20,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x3b,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x53,0x54,0x41,0x47,0x45,0x31,0x5f,0x53,0x49,0x5a,0x45,0x20,0x31,0x34,
+    0x37,0x32,0x35,0x33,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x42,0x49,0x54,0x53,0x20,
+    0x31,0x31,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x53,0x49,0x5a,0x45,0x20,0x28,0x31,
+    0x20,0x3c,0x3c,0x20,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x42,0x49,0x54,0x53,0x29,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,
+    0x46,0x49,0x4e,0x41,0x4c,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x42,0x41,0x54,0x43,0x48,0x5f,0x53,0x49,0x5a,0x45,0x20,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,
+    0x4f,0x52,0x54,0x5f,0x53,0x49,0x5a,0x45,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x49,0x4e,0x41,0x4c,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x4f,0x56,0x45,0x52,
+    0x4c,0x41,0x50,0x5f,0x53,0x49,0x5a,0x45,0x20,0x33,0x32,0x0a,0x5f,0x5f,0x61,0x74,0x74,0x72,0x69,0x62,0x75,0x74,0x65,0x5f,0x5f,0x28,0x28,0x72,0x65,0x71,0x64,0x5f,
+    0x77,0x6f,0x72,0x6b,0x5f,0x67,0x72,0x6f,0x75,0x70,0x5f,0x73,0x69,0x7a,0x65,0x28,0x42,0x57,0x54,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x2c,0x31,
+    0x2c,0x31,0x29,0x29,0x29,0x0a,0x5f,0x5f,0x6b,0x65,0x72,0x6e,0x65,0x6c,0x20,0x76,0x6f,0x69,0x64,0x20,0x42,0x57,0x54,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,
+    0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x20,0x64,0x61,0x74,0x61,0x73,0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,
+    0x74,0x2a,0x20,0x64,0x61,0x74,0x61,0x5f,0x73,0x69,0x7a,0x65,0x73,0x2c,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x64,0x61,0x74,0x61,0x5f,0x73,0x74,0x72,0x69,
+    0x64,0x65,0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x2c,0x5f,0x5f,
+    0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x74,0x6d,0x70,0x5f,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x29,0x0a,0x7b,0x0a,
+    0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x74,0x69,0x64,0x3d,0x67,0x65,0x74,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x5f,0x69,0x64,0x28,
+    0x30,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x69,0x64,0x3d,0x67,0x65,0x74,0x5f,0x67,0x72,0x6f,0x75,0x70,
+    0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x69,0x6e,0x74,0x20,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5b,0x43,0x4f,0x55,
+    0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x53,0x49,0x5a,0x45,0x5d,0x5b,0x32,0x5d,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,
+    0x5f,0x74,0x20,0x69,0x3d,0x74,0x69,0x64,0x3b,0x20,0x69,0x3c,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x53,0x49,0x5a,0x45,0x2a,0x32,
+    0x3b,0x20,0x69,0x2b,0x3d,0x42,0x57,0x54,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x28,0x28,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x69,
+    0x6e,0x74,0x2a,0x29,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x29,0x5b,0x69,0x5d,0x3d,0x30,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,
+    0x5f,0x74,0x20,0x64,0x61,0x74,0x61,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x3d,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x28,0x67,0x69,0x64,0x29,0x2a,0x64,
+    0x61,0x74,0x61,0x5f,0x73,0x74,0x72,0x69,0x64,0x65,0x3b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x20,0x69,0x6e,
+    0x70,0x75,0x74,0x3d,0x64,0x61,0x74,0x61,0x73,0x2b,0x64,0x61,0x74,0x61,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x2b,0x31,0x32,0x38,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,
+    0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x4e,0x3d,0x64,0x61,0x74,0x61,0x5f,0x73,0x69,0x7a,0x65,0x73,0x5b,0x67,0x69,0x64,0x5d,0x2b,0x31,0x3b,0x0a,0x5f,
+    0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x70,0x3d,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,
+    0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x29,0x28,0x69,0x6e,0x70,0x75,0x74,0x29,0x3b,0x0a,0x76,0x6f,0x6c,0x61,0x74,0x69,0x6c,0x65,0x20,0x5f,0x5f,0x6c,0x6f,0x63,0x61,
+    0x6c,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x20,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5f,0x61,0x74,0x6f,0x6d,0x69,0x63,0x3d,0x28,0x76,0x6f,0x6c,0x61,
+    0x74,0x69,0x6c,0x65,0x20,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x29,0x28,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x29,
+    0x3b,0x0a,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x2b,0x3d,0x64,0x61,0x74,0x61,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x3b,0x0a,0x74,0x6d,0x70,0x5f,0x69,0x6e,0x64,0x69,
+    0x63,0x65,0x73,0x2b,0x3d,0x64,0x61,0x74,0x61,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,
+    0x69,0x3d,0x74,0x69,0x64,0x3b,0x20,0x69,0x3c,0x4e,0x3b,0x20,0x69,0x2b,0x3d,0x42,0x57,0x54,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x7b,
+    0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x6e,0x64,0x65,0x78,0x3d,0x69,0x3e,0x3e,0x33,0x3b,0x0a,0x63,0x6f,0x6e,0x73,
+    0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x62,0x69,0x74,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x3d,0x28,0x69,0x26,0x37,0x29,0x3c,0x3c,0x33,0x3b,0x0a,
+    0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x61,0x3d,0x70,0x5b,0x69,0x6e,0x64,0x65,0x78,0x5d,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x36,
+    0x34,0x5f,0x74,0x20,0x62,0x3d,0x70,0x5b,0x69,0x6e,0x64,0x65,0x78,0x2b,0x31,0x5d,0x3b,0x0a,0x69,0x66,0x28,0x62,0x69,0x74,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x3d,
+    0x3d,0x30,0x29,0x0a,0x62,0x3d,0x30,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x76,0x61,0x6c,0x75,0x65,0x3d,0x28,0x61,0x3e,0x3e,0x62,0x69,0x74,0x5f,
+    0x6f,0x66,0x66,0x73,0x65,0x74,0x29,0x7c,0x28,0x62,0x3c,0x3c,0x28,0x36,0x34,0x2d,0x62,0x69,0x74,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x29,0x29,0x3b,0x0a,0x75,0x69,
+    0x6e,0x74,0x32,0x20,0x74,0x6d,0x70,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x63,0x68,0x61,0x72,0x34,0x20,0x6d,0x61,0x73,0x6b,0x3d,0x28,0x75,0x63,0x68,0x61,
+    0x72,0x34,0x29,0x28,0x33,0x2c,0x32,0x2c,0x31,0x2c,0x30,0x29,0x3b,0x0a,0x74,0x6d,0x70,0x2e,0x78,0x3d,0x61,0x73,0x5f,0x75,0x69,0x6e,0x74,0x28,0x73,0x68,0x75,0x66,
+    0x66,0x6c,0x65,0x28,0x61,0x73,0x5f,0x75,0x63,0x68,0x61,0x72,0x34,0x28,0x61,0x73,0x5f,0x75,0x69,0x6e,0x74,0x32,0x28,0x76,0x61,0x6c,0x75,0x65,0x29,0x2e,0x79,0x29,
+    0x2c,0x6d,0x61,0x73,0x6b,0x29,0x29,0x3b,0x0a,0x74,0x6d,0x70,0x2e,0x79,0x3d,0x61,0x73,0x5f,0x75,0x69,0x6e,0x74,0x28,0x73,0x68,0x75,0x66,0x66,0x6c,0x65,0x28,0x61,
+    0x73,0x5f,0x75,0x63,0x68,0x61,0x72,0x34,0x28,0x61,0x73,0x5f,0x75,0x69,0x6e,0x74,0x32,0x28,0x76,0x61,0x6c,0x75,0x65,0x29,0x2e,0x78,0x29,0x2c,0x6d,0x61,0x73,0x6b,
+    0x29,0x29,0x3b,0x0a,0x76,0x61,0x6c,0x75,0x65,0x3d,0x61,0x73,0x5f,0x75,0x6c,0x6f,0x6e,0x67,0x28,0x74,0x6d,0x70,0x29,0x3b,0x0a,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,
+    0x5b,0x69,0x5d,0x3d,0x28,0x76,0x61,0x6c,0x75,0x65,0x26,0x28,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x28,0x2d,0x31,0x29,0x3c,0x3c,0x32,0x31,0x29,0x29,
+    0x7c,0x69,0x3b,0x0a,0x61,0x74,0x6f,0x6d,0x69,0x63,0x5f,0x61,0x64,0x64,0x28,0x28,0x76,0x6f,0x6c,0x61,0x74,0x69,0x6c,0x65,0x20,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,
+    0x20,0x69,0x6e,0x74,0x2a,0x29,0x28,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5f,0x61,0x74,0x6f,0x6d,0x69,0x63,0x2b,0x28,0x28,0x28,0x76,0x61,0x6c,0x75,0x65,0x3e,
+    0x3e,0x28,0x36,0x34,0x2d,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x42,0x49,0x54,0x53,0x2a,0x32,0x29,0x29,0x26,0x28,0x43,0x4f,0x55,
+    0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x53,0x49,0x5a,0x45,0x2d,0x31,0x29,0x29,0x3c,0x3c,0x33,0x29,0x29,0x2c,0x31,0x29,0x3b,0x0a,0x61,0x74,0x6f,
+    0x6d,0x69,0x63,0x5f,0x61,0x64,0x64,0x28,0x28,0x76,0x6f,0x6c,0x61,0x74,0x69,0x6c,0x65,0x20,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x69,0x6e,0x74,0x2a,0x29,0x28,
+    0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5f,0x61,0x74,0x6f,0x6d,0x69,0x63,0x2b,0x28,0x28,0x76,0x61,0x6c,0x75,0x65,0x3e,0x3e,0x28,0x36,0x34,0x2d,0x43,0x4f,0x55,
+    0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x42,0x49,0x54,0x53,0x29,0x29,0x3c,0x3c,0x33,0x29,0x2b,0x34,0x29,0x2c,0x31,0x29,0x3b,0x0a,0x7d,0x0a,0x69,
+    0x66,0x28,0x74,0x69,0x64,0x3d,0x3d,0x30,0x29,0x0a,0x7b,0x0a,0x69,0x6e,0x74,0x20,0x74,0x30,0x3d,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5b,0x30,0x5d,0x5b,0x30,
+    0x5d,0x3b,0x0a,0x69,0x6e,0x74,0x20,0x74,0x31,0x3d,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5b,0x30,0x5d,0x5b,0x31,0x5d,0x3b,0x0a,0x63,0x6f,0x75,0x6e,0x74,0x65,
+    0x72,0x73,0x5b,0x30,0x5d,0x5b,0x30,0x5d,0x3d,0x74,0x30,0x2d,0x31,0x3b,0x0a,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5b,0x30,0x5d,0x5b,0x31,0x5d,0x3d,0x74,0x31,
+    0x2d,0x31,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x3d,0x31,0x3b,0x20,0x69,0x3c,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,
+    0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x53,0x49,0x5a,0x45,0x3b,0x20,0x2b,0x2b,0x69,0x29,0x0a,0x7b,0x0a,0x74,0x30,0x2b,0x3d,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,
+    0x5b,0x69,0x5d,0x5b,0x30,0x5d,0x3b,0x0a,0x74,0x31,0x2b,0x3d,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5b,0x69,0x5d,0x5b,0x31,0x5d,0x3b,0x0a,0x63,0x6f,0x75,0x6e,
+    0x74,0x65,0x72,0x73,0x5b,0x69,0x5d,0x5b,0x30,0x5d,0x3d,0x74,0x30,0x2d,0x31,0x3b,0x0a,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5b,0x69,0x5d,0x5b,0x31,0x5d,0x3d,
+    0x74,0x31,0x2d,0x31,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x74,0x69,0x64,0x3b,0x20,0x69,0x3c,0x4e,0x3b,0x20,0x69,
+    0x2b,0x3d,0x42,0x57,0x54,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,
+    0x5f,0x74,0x20,0x64,0x61,0x74,0x61,0x3d,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x5b,0x69,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x69,0x6e,0x74,0x20,0x6b,0x3d,
+    0x61,0x74,0x6f,0x6d,0x69,0x63,0x5f,0x73,0x75,0x62,0x28,0x28,0x76,0x6f,0x6c,0x61,0x74,0x69,0x6c,0x65,0x20,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x69,0x6e,0x74,
+    0x2a,0x29,0x28,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5f,0x61,0x74,0x6f,0x6d,0x69,0x63,0x2b,0x28,0x28,0x28,0x64,0x61,0x74,0x61,0x3e,0x3e,0x28,0x36,0x34,0x2d,
+    0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x42,0x49,0x54,0x53,0x2a,0x32,0x29,0x29,0x26,0x28,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,
+    0x5f,0x53,0x4f,0x52,0x54,0x5f,0x53,0x49,0x5a,0x45,0x2d,0x31,0x29,0x29,0x3c,0x3c,0x33,0x29,0x29,0x2c,0x31,0x29,0x3b,0x0a,0x74,0x6d,0x70,0x5f,0x69,0x6e,0x64,0x69,
+    0x63,0x65,0x73,0x5b,0x6b,0x5d,0x3d,0x64,0x61,0x74,0x61,0x3b,0x0a,0x7d,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x4e,0x2d,0x31,0x2d,0x74,0x69,
+    0x64,0x3b,0x20,0x69,0x3e,0x3d,0x30,0x3b,0x20,0x69,0x2d,0x3d,0x42,0x57,0x54,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x7b,0x0a,0x63,0x6f,
+    0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x64,0x61,0x74,0x61,0x3d,0x74,0x6d,0x70,0x5f,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x5b,0x69,0x5d,
+    0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x69,0x6e,0x74,0x20,0x6b,0x3d,0x61,0x74,0x6f,0x6d,0x69,0x63,0x5f,0x73,0x75,0x62,0x28,0x28,0x76,0x6f,0x6c,0x61,0x74,0x69,
+    0x6c,0x65,0x20,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x69,0x6e,0x74,0x2a,0x29,0x28,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5f,0x61,0x74,0x6f,0x6d,0x69,0x63,
+    0x2b,0x28,0x28,0x64,0x61,0x74,0x61,0x3e,0x3e,0x28,0x36,0x34,0x2d,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x42,0x49,0x54,0x53,0x29,
+    0x29,0x3c,0x3c,0x33,0x29,0x2b,0x34,0x29,0x2c,0x31,0x29,0x3b,0x0a,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x5b,0x6b,0x5d,0x3d,0x64,0x61,0x74,0x61,0x3b,0x0a,0x7d,0x0a,
+    0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x62,0x75,0x66,0x3d,0x28,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,
+    0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x29,0x28,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x29,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,
+    0x5f,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x4e,0x3b,0x20,0x69,0x2b,0x3d,0x46,0x49,0x4e,0x41,0x4c,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x42,0x41,0x54,0x43,0x48,
+    0x5f,0x53,0x49,0x5a,0x45,0x2d,0x46,0x49,0x4e,0x41,0x4c,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x4f,0x56,0x45,0x52,0x4c,0x41,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x7b,
+    0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6c,0x65,0x6e,0x3d,0x28,0x4e,0x2d,0x69,0x3c,0x46,0x49,0x4e,0x41,0x4c,0x5f,0x53,
+    0x4f,0x52,0x54,0x5f,0x42,0x41,0x54,0x43,0x48,0x5f,0x53,0x49,0x5a,0x45,0x29,0x3f,0x28,0x4e,0x2d,0x69,0x29,0x3a,0x46,0x49,0x4e,0x41,0x4c,0x5f,0x53,0x4f,0x52,0x54,
+    0x5f,0x42,0x41,0x54,0x43,0x48,0x5f,0x53,0x49,0x5a,0x45,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x3d,0x74,0x69,0x64,
+    0x3b,0x20,0x6a,0x3c,0x6c,0x65,0x6e,0x3b,0x20,0x6a,0x2b,0x3d,0x42,0x57,0x54,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x62,0x75,0x66,0x5b,
+    0x6a,0x5d,0x3d,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x5b,0x69,0x2b,0x6a,0x5d,0x3b,0x0a,0x69,0x66,0x28,0x74,0x69,0x64,0x3d,0x3d,0x30,0x29,0x0a,0x7b,0x0a,0x75,0x69,
+    0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x70,0x72,0x65,0x76,0x5f,0x74,0x3d,0x62,0x75,0x66,0x5b,0x30,0x5d,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,
+    0x3d,0x31,0x3b,0x20,0x69,0x3c,0x6c,0x65,0x6e,0x3b,0x20,0x2b,0x2b,0x69,0x29,0x0a,0x7b,0x0a,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x74,0x3d,0x62,0x75,0x66,
+    0x5b,0x69,0x5d,0x3b,0x0a,0x69,0x66,0x28,0x74,0x3c,0x70,0x72,0x65,0x76,0x5f,0x74,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,
+    0x5f,0x74,0x20,0x74,0x32,0x3d,0x70,0x72,0x65,0x76,0x5f,0x74,0x3b,0x0a,0x69,0x6e,0x74,0x20,0x6a,0x3d,0x69,0x2d,0x31,0x3b,0x0a,0x64,0x6f,0x0a,0x7b,0x0a,0x62,0x75,
+    0x66,0x5b,0x6a,0x2b,0x31,0x5d,0x3d,0x70,0x72,0x65,0x76,0x5f,0x74,0x3b,0x0a,0x2d,0x2d,0x6a,0x3b,0x0a,0x69,0x66,0x28,0x6a,0x3c,0x30,0x29,0x0a,0x62,0x72,0x65,0x61,
+    0x6b,0x3b,0x0a,0x70,0x72,0x65,0x76,0x5f,0x74,0x3d,0x62,0x75,0x66,0x5b,0x6a,0x5d,0x3b,0x0a,0x7d,0x20,0x77,0x68,0x69,0x6c,0x65,0x20,0x28,0x74,0x3c,0x70,0x72,0x65,
+    0x76,0x5f,0x74,0x29,0x3b,0x0a,0x62,0x75,0x66,0x5b,0x6a,0x2b,0x31,0x5d,0x3d,0x74,0x3b,0x0a,0x74,0x3d,0x74,0x32,0x3b,0x0a,0x7d,0x0a,0x70,0x72,0x65,0x76,0x5f,0x74,
+    0x3d,0x74,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x3d,0x74,0x69,0x64,0x3b,0x20,0x6a,0x3c,0x6c,
+    0x65,0x6e,0x3b,0x20,0x6a,0x2b,0x3d,0x42,0x57,0x54,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x5b,0x69,
+    0x2b,0x6a,0x5d,0x3d,0x62,0x75,0x66,0x5b,0x6a,0x5d,0x3b,0x0a,0x7d,0x0a,0x2d,0x2d,0x69,0x6e,0x70,0x75,0x74,0x3b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,
+    0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x20,0x6f,0x75,0x74,0x70,0x75,0x74,0x3d,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,
+    0x74,0x2a,0x29,0x28,0x74,0x6d,0x70,0x5f,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x29,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x74,0x69,0x64,
+    0x3b,0x20,0x69,0x3c,0x3d,0x4e,0x3b,0x20,0x69,0x2b,0x3d,0x42,0x57,0x54,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x6f,0x75,0x74,0x70,0x75,
+    0x74,0x5b,0x69,0x5d,0x3d,0x69,0x6e,0x70,0x75,0x74,0x5b,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x5b,0x69,0x5d,0x26,0x28,0x28,0x31,0x3c,0x3c,0x32,0x31,0x29,0x2d,0x31,
+    0x29,0x5d,0x3b,0x0a,0x7d,0x0a,0x5f,0x5f,0x6b,0x65,0x72,0x6e,0x65,0x6c,0x20,0x76,0x6f,0x69,0x64,0x20,0x66,0x69,0x6c,0x74,0x65,0x72,0x28,0x75,0x69,0x6e,0x74,0x33,
+    0x32,0x5f,0x74,0x20,0x6e,0x6f,0x6e,0x63,0x65,0x2c,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x62,0x77,0x74,0x5f,0x6d,0x61,0x78,0x5f,0x73,0x69,0x7a,0x65,0x2c,
+    0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x68,0x61,0x73,0x68,0x65,0x73,0x2c,
+    0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,
+    0x65,0x73,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x3d,0x67,
+    0x65,0x74,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,
+    0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x68,0x61,0x73,0x68,0x3d,0x68,0x61,0x73,0x68,0x65,0x73,0x2b,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x2a,
+    0x28,0x33,0x32,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x29,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,
+    0x74,0x33,0x32,0x5f,0x74,0x20,0x73,0x74,0x61,0x67,0x65,0x32,0x5f,0x73,0x69,0x7a,0x65,0x3d,0x53,0x54,0x41,0x47,0x45,0x31,0x5f,0x53,0x49,0x5a,0x45,0x2b,0x28,0x2a,
+    0x68,0x61,0x73,0x68,0x26,0x30,0x78,0x66,0x66,0x66,0x66,0x66,0x29,0x3b,0x0a,0x69,0x66,0x28,0x73,0x74,0x61,0x67,0x65,0x32,0x5f,0x73,0x69,0x7a,0x65,0x3c,0x62,0x77,
+    0x74,0x5f,0x6d,0x61,0x78,0x5f,0x73,0x69,0x7a,0x65,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x69,0x6e,0x74,0x20,0x69,0x6e,0x64,0x65,0x78,0x3d,0x61,0x74,
+    0x6f,0x6d,0x69,0x63,0x5f,0x61,0x64,0x64,0x28,0x28,0x76,0x6f,0x6c,0x61,0x74,0x69,0x6c,0x65,0x20,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x69,0x6e,0x74,0x2a,
+    0x29,0x28,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x29,0x2c,0x31,0x29,0x2a,0x28,0x33,0x36,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,
+    0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x29,0x29,0x2b,0x31,0x3b,0x0a,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,0x69,
+    0x6e,0x64,0x65,0x78,0x5d,0x3d,0x6e,0x6f,0x6e,0x63,0x65,0x2b,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x3b,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,
+    0x6e,0x72,0x6f,0x6c,0x6c,0x28,0x38,0x29,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x38,0x3b,
+    0x20,0x2b,0x2b,0x69,0x29,0x0a,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,0x69,0x6e,0x64,0x65,0x78,0x2b,0x69,0x2b,0x31,0x5d,
+    0x3d,0x68,0x61,0x73,0x68,0x5b,0x69,0x5d,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x5f,0x5f,0x6b,0x65,0x72,0x6e,0x65,0x6c,0x20,0x76,0x6f,0x69,0x64,0x20,0x70,0x72,0x65,0x70,
+    0x61,0x72,0x65,0x5f,0x62,0x61,0x74,0x63,0x68,0x32,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x68,0x61,
+    0x73,0x68,0x65,0x73,0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,
+    0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x64,0x61,0x74,0x61,0x5f,
+    0x73,0x69,0x7a,0x65,0x73,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,
+    0x64,0x3d,0x67,0x65,0x74,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,
+    0x5f,0x74,0x20,0x4e,0x3d,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,0x30,0x5d,0x2d,0x67,0x65,0x74,0x5f,0x67,0x6c,0x6f,0x62,
+    0x61,0x6c,0x5f,0x73,0x69,0x7a,0x65,0x28,0x30,0x29,0x3b,0x0a,0x69,0x66,0x28,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x3d,0x3d,0x30,0x29,0x0a,0x66,0x69,0x6c,
+    0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,0x30,0x5d,0x3d,0x4e,0x3b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,
+    0x33,0x32,0x5f,0x74,0x2a,0x20,0x68,0x61,0x73,0x68,0x3d,0x68,0x61,0x73,0x68,0x65,0x73,0x2b,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x2a,0x38,0x3b,0x0a,0x5f,
+    0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x3d,
+    0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x2b,0x28,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x2b,0x4e,0x29,0x2a,0x39,0x2b,
+    0x31,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x73,0x74,0x61,0x67,0x65,0x32,0x5f,0x73,0x69,0x7a,0x65,0x3d,0x53,0x54,
+    0x41,0x47,0x45,0x31,0x5f,0x53,0x49,0x5a,0x45,0x2b,0x28,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x5b,0x31,0x5d,0x26,0x30,0x78,0x66,0x66,
+    0x66,0x66,0x66,0x29,0x3b,0x0a,0x64,0x61,0x74,0x61,0x5f,0x73,0x69,0x7a,0x65,0x73,0x5b,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x5d,0x3d,0x73,0x74,0x61,0x67,
+    0x65,0x32,0x5f,0x73,0x69,0x7a,0x65,0x3b,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x28,0x38,0x29,0x0a,0x66,0x6f,0x72,0x20,0x28,
+    0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x38,0x3b,0x20,0x2b,0x2b,0x69,0x29,0x0a,0x68,0x61,0x73,0x68,0x5b,0x69,0x5d,0x3d,
+    0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x5b,0x69,0x2b,0x31,0x5d,0x3b,0x0a,0x7d,0x0a,0x5f,0x5f,0x6b,0x65,0x72,0x6e,0x65,0x6c,0x20,0x76,
+    0x6f,0x69,0x64,0x20,0x66,0x69,0x6e,0x64,0x5f,0x73,0x68,0x61,0x72,0x65,0x73,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,
+    0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x68,0x61,0x73,0x68,0x65,0x73,0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,
+    0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x2c,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,
+    0x74,0x20,0x74,0x61,0x72,0x67,0x65,0x74,0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x73,0x68,0x61,0x72,
+    0x65,0x73,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x6e,0x64,0x65,
+    0x78,0x3d,0x67,0x65,0x74,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x69,0x66,0x28,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,0x67,0x6c,
+    0x6f,0x62,0x61,0x6c,0x5f,0x69,0x6e,0x64,0x65,0x78,0x2a,0x34,0x2b,0x33,0x5d,0x3c,0x74,0x61,0x72,0x67,0x65,0x74,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,
+    0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x64,0x78,0x3d,0x61,0x74,0x6f,0x6d,0x69,0x63,0x5f,0x69,0x6e,0x63,0x28,0x73,0x68,0x61,0x72,0x65,0x73,0x2b,0x30,
+    0x78,0x46,0x46,0x29,0x3b,0x0a,0x69,0x66,0x28,0x69,0x64,0x78,0x3c,0x30,0x78,0x46,0x46,0x29,0x0a,0x73,0x68,0x61,0x72,0x65,0x73,0x5b,0x69,0x64,0x78,0x5d,0x3d,0x66,
+    0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,0x28,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,
+    0x30,0x5d,0x2b,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x6e,0x64,0x65,0x78,0x29,0x2a,0x39,0x2b,0x31,0x5d,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x23,0x64,0x65,0x66,0x69,
+    0x6e,0x65,0x20,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x76,0x2c,0x63,0x29,0x20,0x28,0x72,0x6f,0x74,0x61,0x74,0x65,0x28,0x76,0x2c,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,
+    0x5f,0x74,0x29,0x63,0x29,0x29,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x58,0x4f,0x52,0x28,0x76,0x2c,0x77,0x29,0x20,0x28,0x28,0x76,0x29,0x20,0x5e,0x20,0x28,
+    0x77,0x29,0x29,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x50,0x4c,0x55,0x53,0x28,0x76,0x2c,0x77,0x29,0x20,0x28,0x28,0x76,0x29,0x20,0x2b,0x20,0x28,0x77,0x29,
+    0x29,0x0a,0x5f,0x5f,0x61,0x74,0x74,0x72,0x69,0x62,0x75,0x74,0x65,0x5f,0x5f,0x28,0x28,0x72,0x65,0x71,0x64,0x5f,0x77,0x6f,0x72,0x6b,0x5f,0x67,0x72,0x6f,0x75,0x70,
+    0x5f,0x73,0x69,0x7a,0x65,0x28,0x53,0x41,0x4c,0x53,0x41,0x32,0x30,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x2c,0x31,0x2c,0x31,0x29,0x29,0x29,0x0a,
+    0x5f,0x5f,0x6b,0x65,0x72,0x6e,0x65,0x6c,0x20,0x76,0x6f,0x69,0x64,0x20,0x53,0x61,0x6c,0x73,0x61,0x32,0x30,0x5f,0x58,0x4f,0x52,0x4b,0x65,0x79,0x53,0x74,0x72,0x65,
+    0x61,0x6d,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x6b,0x65,0x79,0x73,
+    0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x6f,0x75,0x74,0x70,0x75,0x74,0x73,0x2c,0x5f,0x5f,0x67,0x6c,
+    0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x73,0x2c,0x75,0x69,0x6e,0x74,
+    0x33,0x32,0x5f,0x74,0x20,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x74,0x72,0x69,0x64,0x65,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,
+    0x33,0x32,0x5f,0x74,0x20,0x74,0x3d,0x67,0x65,0x74,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,
+    0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x3d,0x67,0x65,0x74,0x5f,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,
+    0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x3d,0x67,0x2a,0x28,0x28,0x75,0x69,0x6e,0x74,0x36,
+    0x34,0x5f,0x74,0x29,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x74,0x72,0x69,0x64,0x65,0x29,0x2b,0x31,0x32,0x38,0x3b,0x0a,0x7b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,
+    0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x70,0x3d,0x6f,0x75,0x74,0x70,0x75,0x74,0x73,0x2b,0x28,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x6f,
+    0x66,0x66,0x73,0x65,0x74,0x2d,0x31,0x32,0x38,0x29,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x29,0x3b,0x0a,0x66,0x6f,0x72,
+    0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x3d,0x74,0x3b,0x20,0x69,0x3c,0x31,0x32,0x38,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,
+    0x74,0x33,0x32,0x5f,0x74,0x29,0x3b,0x20,0x69,0x2b,0x3d,0x53,0x41,0x4c,0x53,0x41,0x32,0x30,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x70,
+    0x5b,0x69,0x5d,0x3d,0x30,0x3b,0x0a,0x7d,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,
+    0x2a,0x20,0x6b,0x3d,0x6b,0x65,0x79,0x73,0x2b,0x67,0x2a,0x38,0x3b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,
+    0x20,0x6f,0x75,0x74,0x70,0x75,0x74,0x3d,0x6f,0x75,0x74,0x70,0x75,0x74,0x73,0x2b,0x28,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x2b,0x28,
+    0x74,0x2a,0x36,0x34,0x29,0x29,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,
+    0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x3d,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x73,
+    0x5b,0x67,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x31,0x3d,0x6b,0x5b,0x30,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,
+    0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x32,0x3d,0x6b,0x5b,0x31,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,
+    0x32,0x5f,0x74,0x20,0x6a,0x33,0x3d,0x6b,0x5b,0x32,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x34,0x3d,0x6b,
+    0x5b,0x33,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x31,0x31,0x3d,0x6b,0x5b,0x34,0x5d,0x3b,0x0a,0x63,0x6f,
+    0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x31,0x32,0x3d,0x6b,0x5b,0x35,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,
+    0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x31,0x33,0x3d,0x6b,0x5b,0x36,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,
+    0x31,0x34,0x3d,0x6b,0x5b,0x37,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x30,0x3d,0x30,0x78,0x36,0x31,0x37,
+    0x30,0x37,0x38,0x36,0x35,0x55,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x35,0x3d,0x30,0x78,0x33,0x33,0x32,0x30,
+    0x36,0x34,0x36,0x45,0x55,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x31,0x30,0x3d,0x30,0x78,0x37,0x39,0x36,0x32,
+    0x32,0x44,0x33,0x32,0x55,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x31,0x35,0x3d,0x30,0x78,0x36,0x42,0x32,0x30,
+    0x36,0x35,0x37,0x34,0x55,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x36,0x3d,0x30,0x3b,0x0a,0x63,0x6f,0x6e,0x73,
+    0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x37,0x3d,0x30,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,
+    0x6a,0x38,0x3d,0x30,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x39,0x3d,0x30,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,
+    0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x3d,0x74,0x2a,0x36,0x34,0x3b,0x20,0x69,0x3c,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x3b,0x20,
+    0x69,0x2b,0x3d,0x53,0x41,0x4c,0x53,0x41,0x32,0x30,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x2a,0x36,0x34,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,
+    0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x38,0x5f,0x31,0x3d,0x6a,0x38,0x2b,0x28,0x69,0x2f,0x36,0x34,0x29,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,
+    0x32,0x5f,0x74,0x20,0x78,0x30,0x3d,0x6a,0x30,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x31,0x3d,0x6a,0x31,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,
+    0x32,0x5f,0x74,0x20,0x78,0x32,0x3d,0x6a,0x32,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x33,0x3d,0x6a,0x33,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,
+    0x32,0x5f,0x74,0x20,0x78,0x34,0x3d,0x6a,0x34,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x35,0x3d,0x6a,0x35,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,
+    0x32,0x5f,0x74,0x20,0x78,0x36,0x3d,0x6a,0x36,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x37,0x3d,0x6a,0x37,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,
+    0x32,0x5f,0x74,0x20,0x78,0x38,0x3d,0x6a,0x38,0x5f,0x31,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x39,0x3d,0x6a,0x39,0x3b,0x0a,0x75,0x69,0x6e,
+    0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x31,0x30,0x3d,0x6a,0x31,0x30,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x31,0x31,0x3d,0x6a,0x31,0x31,0x3b,
+    0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x31,0x32,0x3d,0x6a,0x31,0x32,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x31,0x33,0x3d,
+    0x6a,0x31,0x33,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x31,0x34,0x3d,0x6a,0x31,0x34,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,
+    0x78,0x31,0x35,0x3d,0x6a,0x31,0x35,0x3b,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x28,0x35,0x29,0x0a,0x66,0x6f,0x72,0x20,0x28,
+    0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x3d,0x30,0x3b,0x20,0x6a,0x3c,0x31,0x30,0x3b,0x20,0x2b,0x2b,0x6a,0x29,0x0a,0x7b,0x0a,0x78,0x34,0x3d,0x58,0x4f,
+    0x52,0x28,0x20,0x78,0x34,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x30,0x2c,0x78,0x31,0x32,0x29,0x2c,0x37,0x29,0x29,0x3b,0x0a,
+    0x78,0x38,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x38,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x34,0x2c,0x78,0x30,0x29,0x2c,0x39,
+    0x29,0x29,0x3b,0x0a,0x78,0x31,0x32,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x32,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x38,0x2c,
+    0x78,0x34,0x29,0x2c,0x31,0x33,0x29,0x29,0x3b,0x0a,0x78,0x30,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x30,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,
+    0x28,0x78,0x31,0x32,0x2c,0x78,0x38,0x29,0x2c,0x31,0x38,0x29,0x29,0x3b,0x0a,0x78,0x39,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x39,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,
+    0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x35,0x2c,0x78,0x31,0x29,0x2c,0x37,0x29,0x29,0x3b,0x0a,0x78,0x31,0x33,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x33,0x2c,0x52,
+    0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x39,0x2c,0x78,0x35,0x29,0x2c,0x39,0x29,0x29,0x3b,0x0a,0x78,0x31,0x3d,0x58,0x4f,0x52,0x28,0x20,
+    0x78,0x31,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x33,0x2c,0x78,0x39,0x29,0x2c,0x31,0x33,0x29,0x29,0x3b,0x0a,0x78,0x35,0x3d,
+    0x58,0x4f,0x52,0x28,0x20,0x78,0x35,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x31,0x2c,0x78,0x31,0x33,0x29,0x2c,0x31,0x38,0x29,
+    0x29,0x3b,0x0a,0x78,0x31,0x34,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x34,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x30,0x2c,0x78,
+    0x36,0x29,0x2c,0x37,0x29,0x29,0x3b,0x0a,0x78,0x32,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x32,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,
+    0x31,0x34,0x2c,0x78,0x31,0x30,0x29,0x2c,0x39,0x29,0x29,0x3b,0x0a,0x78,0x36,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x36,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,
+    0x4c,0x55,0x53,0x28,0x20,0x78,0x32,0x2c,0x78,0x31,0x34,0x29,0x2c,0x31,0x33,0x29,0x29,0x3b,0x0a,0x78,0x31,0x30,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x30,0x2c,0x52,
+    0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x36,0x2c,0x78,0x32,0x29,0x2c,0x31,0x38,0x29,0x29,0x3b,0x0a,0x78,0x33,0x3d,0x58,0x4f,0x52,0x28,
+    0x20,0x78,0x33,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x35,0x2c,0x78,0x31,0x31,0x29,0x2c,0x37,0x29,0x29,0x3b,0x0a,0x78,0x37,
+    0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x37,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x33,0x2c,0x78,0x31,0x35,0x29,0x2c,0x39,0x29,
+    0x29,0x3b,0x0a,0x78,0x31,0x31,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x31,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x37,0x2c,0x78,
+    0x33,0x29,0x2c,0x31,0x33,0x29,0x29,0x3b,0x0a,0x78,0x31,0x35,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x35,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,
+    0x28,0x78,0x31,0x31,0x2c,0x78,0x37,0x29,0x2c,0x31,0x38,0x29,0x29,0x3b,0x0a,0x78,0x31,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x31,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,
+    0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x30,0x2c,0x78,0x33,0x29,0x2c,0x37,0x29,0x29,0x3b,0x0a,0x78,0x32,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x32,0x2c,0x52,0x4f,
+    0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x31,0x2c,0x78,0x30,0x29,0x2c,0x39,0x29,0x29,0x3b,0x0a,0x78,0x33,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,
+    0x33,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x32,0x2c,0x78,0x31,0x29,0x2c,0x31,0x33,0x29,0x29,0x3b,0x0a,0x78,0x30,0x3d,0x58,
+    0x4f,0x52,0x28,0x20,0x78,0x30,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x33,0x2c,0x78,0x32,0x29,0x2c,0x31,0x38,0x29,0x29,0x3b,
+    0x0a,0x78,0x36,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x36,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x35,0x2c,0x78,0x34,0x29,0x2c,
+    0x37,0x29,0x29,0x3b,0x0a,0x78,0x37,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x37,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x36,0x2c,
+    0x78,0x35,0x29,0x2c,0x39,0x29,0x29,0x3b,0x0a,0x78,0x34,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x34,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,
+    0x20,0x78,0x37,0x2c,0x78,0x36,0x29,0x2c,0x31,0x33,0x29,0x29,0x3b,0x0a,0x78,0x35,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x35,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,
+    0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x34,0x2c,0x78,0x37,0x29,0x2c,0x31,0x38,0x29,0x29,0x3b,0x0a,0x78,0x31,0x31,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x31,0x2c,0x52,
+    0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x30,0x2c,0x78,0x39,0x29,0x2c,0x37,0x29,0x29,0x3b,0x0a,0x78,0x38,0x3d,0x58,0x4f,0x52,0x28,0x20,
+    0x78,0x38,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x31,0x2c,0x78,0x31,0x30,0x29,0x2c,0x39,0x29,0x29,0x3b,0x0a,0x78,0x39,0x3d,
+    0x58,0x4f,0x52,0x28,0x20,0x78,0x39,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x38,0x2c,0x78,0x31,0x31,0x29,0x2c,0x31,0x33,0x29,
+    0x29,0x3b,0x0a,0x78,0x31,0x30,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x30,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x39,0x2c,0x78,
+    0x38,0x29,0x2c,0x31,0x38,0x29,0x29,0x3b,0x0a,0x78,0x31,0x32,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x32,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,
+    0x28,0x78,0x31,0x35,0x2c,0x78,0x31,0x34,0x29,0x2c,0x37,0x29,0x29,0x3b,0x0a,0x78,0x31,0x33,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x33,0x2c,0x52,0x4f,0x54,0x41,0x54,
+    0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x32,0x2c,0x78,0x31,0x35,0x29,0x2c,0x39,0x29,0x29,0x3b,0x0a,0x78,0x31,0x34,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x34,
+    0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x33,0x2c,0x78,0x31,0x32,0x29,0x2c,0x31,0x33,0x29,0x29,0x3b,0x0a,0x78,0x31,0x35,0x3d,
+    0x58,0x4f,0x52,0x28,0x78,0x31,0x35,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x34,0x2c,0x78,0x31,0x33,0x29,0x2c,0x31,0x38,0x29,
+    0x29,0x3b,0x0a,0x7d,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x30,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x30,0x2c,0x6a,0x30,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,
+    0x75,0x74,0x5b,0x31,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x2c,0x6a,0x31,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x32,0x5d,0x3d,0x50,0x4c,0x55,
+    0x53,0x28,0x78,0x32,0x2c,0x6a,0x32,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x33,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x33,0x2c,0x6a,0x33,0x29,0x3b,
+    0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x34,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x34,0x2c,0x6a,0x34,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x35,
+    0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x35,0x2c,0x6a,0x35,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x36,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x36,
+    0x2c,0x6a,0x36,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x37,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x37,0x2c,0x6a,0x37,0x29,0x3b,0x0a,0x6f,0x75,0x74,
+    0x70,0x75,0x74,0x5b,0x38,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x38,0x2c,0x6a,0x38,0x5f,0x31,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x39,0x5d,0x3d,
+    0x50,0x4c,0x55,0x53,0x28,0x78,0x39,0x2c,0x6a,0x39,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x31,0x30,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x30,
+    0x2c,0x6a,0x31,0x30,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x31,0x31,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x31,0x2c,0x6a,0x31,0x31,0x29,0x3b,
+    0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x31,0x32,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x32,0x2c,0x6a,0x31,0x32,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,
+    0x74,0x5b,0x31,0x33,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x33,0x2c,0x6a,0x31,0x33,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x31,0x34,0x5d,0x3d,
+    0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x34,0x2c,0x6a,0x31,0x34,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x31,0x35,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,
+    0x31,0x35,0x2c,0x6a,0x31,0x35,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x2b,0x3d,0x28,0x53,0x41,0x4c,0x53,0x41,0x32,0x30,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,
+    0x53,0x49,0x5a,0x45,0x2a,0x36,0x34,0x29,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x29,0x3b,0x0a,0x7d,0x0a,0x62,0x61,0x72,
+    0x72,0x69,0x65,0x72,0x28,0x43,0x4c,0x4b,0x5f,0x47,0x4c,0x4f,0x42,0x41,0x4c,0x5f,0x4d,0x45,0x4d,0x5f,0x46,0x45,0x4e,0x43,0x45,0x29,0x3b,0x0a,0x69,0x66,0x28,0x74,
+    0x3c,0x31,0x36,0x29,0x0a,0x7b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x70,0x3d,0x6f,0x75,0x74,0x70,
+    0x75,0x74,0x73,0x2b,0x28,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x2b,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x2b,0x33,
+    0x29,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x29,0x3b,0x0a,0x70,0x5b,0x74,0x5d,0x3d,0x30,0x3b,0x0a,0x7d,0x0a,0x69,0x66,
+    0x28,0x28,0x74,0x3d,0x3d,0x30,0x29,0x26,0x26,0x28,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x26,0x33,0x29,0x29,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,
+    0x73,0x5b,0x28,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x2b,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x29,0x2f,0x73,0x69,
+    0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x29,0x5d,0x20,0x26,0x3d,0x20,0x30,0x78,0x46,0x46,0x46,0x46,0x46,0x46,0x46,0x46,0x55,0x3e,0x3e,
+    0x28,0x28,0x34,0x2d,0x28,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x26,0x33,0x29,0x29,0x3c,0x3c,0x33,0x29,0x3b,0x0a,0x7d,0x0a,0x23,0x64,0x65,0x66,
+    0x69,0x6e,0x65,0x20,0x52,0x4f,0x55,0x4e,0x44,0x53,0x20,0x32,0x34,0x20,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x52,0x36,0x34,0x28,0x61,0x2c,0x62,0x2c,0x63,
+    0x29,0x20,0x28,0x28,0x28,0x61,0x29,0x20,0x3c,0x3c,0x20,0x62,0x29,0x20,0x7c,0x20,0x28,0x28,0x61,0x29,0x20,0x3e,0x3e,0x20,0x63,0x29,0x29,0x0a,0x5f,0x5f,0x63,0x6f,
+    0x6e,0x73,0x74,0x61,0x6e,0x74,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x72,0x63,0x5b,0x32,0x5d,0x5b,0x52,0x4f,0x55,0x4e,
+    0x44,0x53,0x5d,0x3d,0x7b,0x0a,0x7b,0x30,0x78,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x31,0x55,0x4c,0x2c,0x30,0x78,0x30,0x30,
+    0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x38,0x32,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,
+    0x30,0x38,0x41,0x55,0x4c,0x2c,0x0a,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x55,0x4c,0x2c,0x30,0x78,0x30,0x30,
+    0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x38,0x42,0x55,0x4c,0x2c,0x30,0x78,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x30,
+    0x30,0x30,0x31,0x55,0x4c,0x2c,0x0a,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x38,0x30,0x38,0x31,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,
+    0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x39,0x55,0x4c,0x2c,0x30,0x78,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,
+    0x30,0x38,0x41,0x55,0x4c,0x2c,0x0a,0x30,0x78,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x38,0x55,0x4c,0x2c,0x30,0x78,0x30,0x30,
+    0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x38,0x30,0x30,0x39,0x55,0x4c,0x2c,0x30,0x78,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x30,
+    0x30,0x30,0x41,0x55,0x4c,0x2c,0x0a,0x30,0x78,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x38,0x30,0x38,0x42,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,
+    0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x42,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,
+    0x30,0x38,0x39,0x55,0x4c,0x2c,0x0a,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x33,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,
+    0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x32,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,
+    0x30,0x38,0x30,0x55,0x4c,0x2c,0x0a,0x30,0x78,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x41,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,
+    0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x41,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x38,
+    0x30,0x38,0x31,0x55,0x4c,0x2c,0x0a,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x38,0x30,0x55,0x4c,0x2c,0x30,0x78,0x30,0x30,
+    0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x31,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x38,
+    0x30,0x30,0x38,0x55,0x4c,0x7d,0x2c,0x0a,0x7b,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,
+    0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x0a,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,
+    0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x0a,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,
+    0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x7d,0x0a,0x7d,0x3b,0x0a,0x5f,0x5f,0x63,0x6f,0x6e,0x73,0x74,0x61,0x6e,0x74,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,
+    0x69,0x6e,0x74,0x20,0x72,0x6f,0x5b,0x32,0x35,0x5d,0x5b,0x32,0x5d,0x3d,0x7b,0x0a,0x7b,0x20,0x30,0x2c,0x36,0x34,0x7d,0x2c,0x7b,0x34,0x34,0x2c,0x32,0x30,0x7d,0x2c,
+    0x7b,0x34,0x33,0x2c,0x32,0x31,0x7d,0x2c,0x7b,0x32,0x31,0x2c,0x34,0x33,0x7d,0x2c,0x7b,0x31,0x34,0x2c,0x35,0x30,0x7d,0x2c,0x0a,0x7b,0x20,0x31,0x2c,0x36,0x33,0x7d,
+    0x2c,0x7b,0x20,0x36,0x2c,0x35,0x38,0x7d,0x2c,0x7b,0x32,0x35,0x2c,0x33,0x39,0x7d,0x2c,0x7b,0x20,0x38,0x2c,0x35,0x36,0x7d,0x2c,0x7b,0x31,0x38,0x2c,0x34,0x36,0x7d,
+    0x2c,0x0a,0x7b,0x36,0x32,0x2c,0x32,0x7d,0x2c,0x7b,0x35,0x35,0x2c,0x39,0x7d,0x2c,0x7b,0x33,0x39,0x2c,0x32,0x35,0x7d,0x2c,0x7b,0x34,0x31,0x2c,0x32,0x33,0x7d,0x2c,
+    0x7b,0x20,0x32,0x2c,0x36,0x32,0x7d,0x2c,0x0a,0x7b,0x32,0x38,0x2c,0x33,0x36,0x7d,0x2c,0x7b,0x32,0x30,0x2c,0x34,0x34,0x7d,0x2c,0x7b,0x20,0x33,0x2c,0x36,0x31,0x7d,
+    0x2c,0x7b,0x34,0x35,0x2c,0x31,0x39,0x7d,0x2c,0x7b,0x36,0x31,0x2c,0x33,0x7d,0x2c,0x0a,0x7b,0x32,0x37,0x2c,0x33,0x37,0x7d,0x2c,0x7b,0x33,0x36,0x2c,0x32,0x38,0x7d,
+    0x2c,0x7b,0x31,0x30,0x2c,0x35,0x34,0x7d,0x2c,0x7b,0x31,0x35,0x2c,0x34,0x39,0x7d,0x2c,0x7b,0x35,0x36,0x2c,0x38,0x7d,0x0a,0x7d,0x3b,0x0a,0x5f,0x5f,0x63,0x6f,0x6e,
+    0x73,0x74,0x61,0x6e,0x74,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x69,0x6e,0x74,0x20,0x61,0x5b,0x32,0x35,0x5d,0x3d,0x7b,0x0a,0x30,0x2c,0x36,0x2c,0x31,0x32,0x2c,0x31,
+    0x38,0x2c,0x32,0x34,0x2c,0x0a,0x31,0x2c,0x37,0x2c,0x31,0x33,0x2c,0x31,0x39,0x2c,0x32,0x30,0x2c,0x0a,0x32,0x2c,0x38,0x2c,0x31,0x34,0x2c,0x31,0x35,0x2c,0x32,0x31,
+    0x2c,0x0a,0x33,0x2c,0x39,0x2c,0x31,0x30,0x2c,0x31,0x36,0x2c,0x32,0x32,0x2c,0x0a,0x34,0x2c,0x35,0x2c,0x31,0x31,0x2c,0x31,0x37,0x2c,0x32,0x33,0x0a,0x7d,0x3b,0x0a,
+    0x5f,0x5f,0x63,0x6f,0x6e,0x73,0x74,0x61,0x6e,0x74,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x69,0x6e,0x74,0x20,0x62,0x5b,0x32,0x35,0x5d,0x3d,0x7b,0x0a,0x30,0x2c,0x31,
+    0x2c,0x32,0x2c,0x33,0x2c,0x34,0x2c,0x0a,0x31,0x2c,0x32,0x2c,0x33,0x2c,0x34,0x2c,0x30,0x2c,0x0a,0x32,0x2c,0x33,0x2c,0x34,0x2c,0x30,0x2c,0x31,0x2c,0x0a,0x33,0x2c,
+    0x34,0x2c,0x30,0x2c,0x31,0x2c,0x32,0x2c,0x0a,0x34,0x2c,0x30,0x2c,0x31,0x2c,0x32,0x2c,0x33,0x0a,0x7d,0x3b,0x0a,0x5f,0x5f,0x63,0x6f,0x6e,0x73,0x74,0x61,0x6e,0x74,
+    0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x69,0x6e,0x74,0x20,0x63,0x5b,0x32,0x35,0x5d,0x5b,0x33,0x5d,0x3d,0x7b,0x0a,0x7b,0x20,0x30,0x2c,0x31,0x2c,0x32,0x7d,0x2c,0x7b,
+    0x20,0x31,0x2c,0x32,0x2c,0x33,0x7d,0x2c,0x7b,0x20,0x32,0x2c,0x33,0x2c,0x34,0x7d,0x2c,0x7b,0x20,0x33,0x2c,0x34,0x2c,0x30,0x7d,0x2c,0x7b,0x20,0x34,0x2c,0x30,0x2c,
+    0x31,0x7d,0x2c,0x0a,0x7b,0x20,0x35,0x2c,0x36,0x2c,0x37,0x7d,0x2c,0x7b,0x20,0x36,0x2c,0x37,0x2c,0x38,0x7d,0x2c,0x7b,0x20,0x37,0x2c,0x38,0x2c,0x39,0x7d,0x2c,0x7b,
+    0x20,0x38,0x2c,0x39,0x2c,0x35,0x7d,0x2c,0x7b,0x20,0x39,0x2c,0x35,0x2c,0x36,0x7d,0x2c,0x0a,0x7b,0x31,0x30,0x2c,0x31,0x31,0x2c,0x31,0x32,0x7d,0x2c,0x7b,0x31,0x31,
+    0x2c,0x31,0x32,0x2c,0x31,0x33,0x7d,0x2c,0x7b,0x31,0x32,0x2c,0x31,0x33,0x2c,0x31,0x34,0x7d,0x2c,0x7b,0x31,0x33,0x2c,0x31,0x34,0x2c,0x31,0x30,0x7d,0x2c,0x7b,0x31,
+    0x34,0x2c,0x31,0x30,0x2c,0x31,0x31,0x7d,0x2c,0x0a,0x7b,0x31,0x35,0x2c,0x31,0x36,0x2c,0x31,0x37,0x7d,0x2c,0x7b,0x31,0x36,0x2c,0x31,0x37,0x2c,0x31,0x38,0x7d,0x2c,
+    0x7b,0x31,0x37,0x2c,0x31,0x38,0x2c,0x31,0x39,0x7d,0x2c,0x7b,0x31,0x38,0x2c,0x31,0x39,0x2c,0x31,0x35,0x7d,0x2c,0x7b,0x31,0x39,0x2c,0x31,0x35,0x2c,0x31,0x36,0x7d,
+    0x2c,0x0a,0x7b,0x32,0x30,0x2c,0x32,0x31,0x2c,0x32,0x32,0x7d,0x2c,0x7b,0x32,0x31,0x2c,0x32,0x32,0x2c,0x32,0x33,0x7d,0x2c,0x7b,0x32,0x32,0x2c,0x32,0x33,0x2c,0x32,
+    0x34,0x7d,0x2c,0x7b,0x32,0x33,0x2c,0x32,0x34,0x2c,0x32,0x30,0x7d,0x2c,0x7b,0x32,0x34,0x2c,0x32,0x30,0x2c,0x32,0x31,0x7d,0x0a,0x7d,0x3b,0x0a,0x5f,0x5f,0x63,0x6f,
+    0x6e,0x73,0x74,0x61,0x6e,0x74,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x69,0x6e,0x74,0x20,0x64,0x5b,0x32,0x35,0x5d,0x3d,0x7b,0x0a,0x30,0x2c,0x31,0x2c,0x32,0x2c,0x33,
+    0x2c,0x34,0x2c,0x0a,0x31,0x30,0x2c,0x31,0x31,0x2c,0x31,0x32,0x2c,0x31,0x33,0x2c,0x31,0x34,0x2c,0x0a,0x32,0x30,0x2c,0x32,0x31,0x2c,0x32,0x32,0x2c,0x32,0x33,0x2c,
+    0x32,0x34,0x2c,0x0a,0x35,0x2c,0x36,0x2c,0x37,0x2c,0x38,0x2c,0x39,0x2c,0x0a,0x31,0x35,0x2c,0x31,0x36,0x2c,0x31,0x37,0x2c,0x31,0x38,0x2c,0x31,0x39,0x0a,0x7d,0x3b,
+    0x0a,0x5f,0x5f,0x61,0x74,0x74,0x72,0x69,0x62,0x75,0x74,0x65,0x5f,0x5f,0x28,0x28,0x72,0x65,0x71,0x64,0x5f,0x77,0x6f,0x72,0x6b,0x5f,0x67,0x72,0x6f,0x75,0x70,0x5f,
+    0x73,0x69,0x7a,0x65,0x28,0x33,0x32,0x2c,0x31,0x2c,0x31,0x29,0x29,0x29,0x0a,0x5f,0x5f,0x6b,0x65,0x72,0x6e,0x65,0x6c,0x20,0x76,0x6f,0x69,0x64,0x20,0x73,0x68,0x61,
+    0x33,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x20,0x69,0x6e,0x70,0x75,0x74,0x73,
+    0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x69,0x6e,0x70,0x75,0x74,0x5f,
+    0x73,0x69,0x7a,0x65,0x73,0x2c,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x6e,0x70,0x75,0x74,0x5f,0x73,0x74,0x72,0x69,0x64,0x65,0x2c,0x5f,0x5f,0x67,0x6c,
+    0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x68,0x61,0x73,0x68,0x65,0x73,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,
+    0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x74,0x3d,0x67,0x65,0x74,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,
+    0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x3d,0x67,0x65,0x74,0x5f,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x69,0x66,0x28,
+    0x74,0x3e,0x3d,0x32,0x35,0x29,0x0a,0x72,0x65,0x74,0x75,0x72,0x6e,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x73,0x3d,
+    0x74,0x20,0x25,0x20,0x35,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x69,0x6e,0x70,0x75,0x74,0x5f,0x6f,0x66,0x66,0x73,
+    0x65,0x74,0x3d,0x28,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x69,0x6e,0x70,0x75,0x74,0x5f,0x73,0x74,0x72,0x69,0x64,0x65,0x29,0x2a,0x67,0x3b,0x0a,0x5f,
+    0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x69,0x6e,0x70,0x75,0x74,0x3d,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,
+    0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x29,0x28,0x69,0x6e,0x70,0x75,0x74,0x73,0x2b,0x69,0x6e,0x70,0x75,0x74,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,
+    0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x6e,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x3d,0x69,0x6e,0x70,
+    0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x73,0x5b,0x67,0x5d,0x2b,0x31,0x3b,0x0a,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,
+    0x41,0x5b,0x32,0x35,0x5d,0x3b,0x0a,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x43,0x5b,0x32,0x35,0x5d,0x3b,0x0a,0x5f,
+    0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x44,0x5b,0x32,0x35,0x5d,0x3b,0x0a,0x41,0x5b,0x74,0x5d,0x3d,0x30,0x3b,0x0a,0x63,
+    0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x77,0x6f,0x72,0x64,0x73,0x3d,0x69,0x6e,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x2f,0x73,
+    0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,
+    0x74,0x61,0x69,0x6c,0x5f,0x73,0x69,0x7a,0x65,0x3d,0x69,0x6e,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x20,0x25,0x20,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,
+    0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x77,0x6f,0x72,0x64,0x49,0x6e,0x64,0x65,0x78,0x3d,0x30,0x3b,0x0a,0x66,
+    0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x77,0x6f,0x72,0x64,0x73,0x3b,0x20,0x2b,0x2b,0x69,0x2c,0x2b,
+    0x2b,0x69,0x6e,0x70,0x75,0x74,0x29,0x0a,0x7b,0x0a,0x41,0x5b,0x77,0x6f,0x72,0x64,0x49,0x6e,0x64,0x65,0x78,0x5d,0x20,0x5e,0x3d,0x20,0x2a,0x69,0x6e,0x70,0x75,0x74,
+    0x3b,0x0a,0x2b,0x2b,0x77,0x6f,0x72,0x64,0x49,0x6e,0x64,0x65,0x78,0x3b,0x0a,0x69,0x66,0x28,0x77,0x6f,0x72,0x64,0x49,0x6e,0x64,0x65,0x78,0x3d,0x3d,0x31,0x37,0x29,
+    0x0a,0x7b,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x28,0x52,0x4f,0x55,0x4e,0x44,0x53,0x29,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,
+    0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x52,0x4f,0x55,0x4e,0x44,0x53,0x3b,0x20,0x2b,0x2b,0x69,0x29,0x0a,0x7b,0x0a,0x43,0x5b,0x74,0x5d,0x3d,0x41,0x5b,
+    0x73,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x35,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x31,0x30,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x31,0x35,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x32,0x30,
+    0x5d,0x3b,0x0a,0x44,0x5b,0x74,0x5d,0x3d,0x43,0x5b,0x62,0x5b,0x32,0x30,0x2b,0x73,0x5d,0x5d,0x5e,0x52,0x36,0x34,0x28,0x43,0x5b,0x62,0x5b,0x35,0x2b,0x73,0x5d,0x5d,
+    0x2c,0x31,0x2c,0x36,0x33,0x29,0x3b,0x0a,0x43,0x5b,0x74,0x5d,0x3d,0x52,0x36,0x34,0x28,0x41,0x5b,0x61,0x5b,0x74,0x5d,0x5d,0x5e,0x44,0x5b,0x62,0x5b,0x74,0x5d,0x5d,
+    0x2c,0x72,0x6f,0x5b,0x74,0x5d,0x5b,0x30,0x5d,0x2c,0x72,0x6f,0x5b,0x74,0x5d,0x5b,0x31,0x5d,0x29,0x3b,0x0a,0x41,0x5b,0x64,0x5b,0x74,0x5d,0x5d,0x3d,0x43,0x5b,0x63,
+    0x5b,0x74,0x5d,0x5b,0x30,0x5d,0x5d,0x5e,0x28,0x28,0x7e,0x43,0x5b,0x63,0x5b,0x74,0x5d,0x5b,0x31,0x5d,0x5d,0x29,0x26,0x43,0x5b,0x63,0x5b,0x74,0x5d,0x5b,0x32,0x5d,
+    0x5d,0x29,0x3b,0x0a,0x41,0x5b,0x74,0x5d,0x20,0x5e,0x3d,0x20,0x72,0x63,0x5b,0x28,0x74,0x3d,0x3d,0x30,0x29,0x3f,0x30,0x3a,0x31,0x5d,0x5b,0x69,0x5d,0x3b,0x20,0x0a,
+    0x7d,0x0a,0x77,0x6f,0x72,0x64,0x49,0x6e,0x64,0x65,0x78,0x3d,0x30,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x74,0x61,0x69,0x6c,
+    0x3d,0x30,0x3b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x20,0x70,0x3d,0x28,0x5f,
+    0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x29,0x69,0x6e,0x70,0x75,0x74,0x3b,0x0a,0x66,0x6f,
+    0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x74,0x61,0x69,0x6c,0x5f,0x73,0x69,0x7a,0x65,0x3b,0x20,0x2b,0x2b,
+    0x69,0x29,0x0a,0x7b,0x0a,0x74,0x61,0x69,0x6c,0x7c,0x3d,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x28,0x70,0x5b,0x69,0x5d,0x29,0x3c,0x3c,0x28,0x69,0x2a,
+    0x38,0x29,0x3b,0x0a,0x7d,0x0a,0x41,0x5b,0x77,0x6f,0x72,0x64,0x49,0x6e,0x64,0x65,0x78,0x5d,0x20,0x5e,0x3d,0x20,0x74,0x61,0x69,0x6c,0x5e,0x28,0x28,0x75,0x69,0x6e,
+    0x74,0x36,0x34,0x5f,0x74,0x29,0x28,0x28,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x28,0x30,0x78,0x30,0x32,0x7c,0x28,0x31,0x3c,0x3c,0x32,0x29,0x29,0x29,
+    0x3c,0x3c,0x28,0x74,0x61,0x69,0x6c,0x5f,0x73,0x69,0x7a,0x65,0x2a,0x38,0x29,0x29,0x29,0x3b,0x0a,0x41,0x5b,0x31,0x36,0x5d,0x20,0x5e,0x3d,0x20,0x30,0x78,0x38,0x30,
+    0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x55,0x4c,0x3b,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,
+    0x28,0x31,0x29,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x52,0x4f,0x55,0x4e,0x44,0x53,0x3b,0x20,0x2b,0x2b,0x69,0x29,
+    0x0a,0x7b,0x0a,0x43,0x5b,0x74,0x5d,0x3d,0x41,0x5b,0x73,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x35,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x31,0x30,0x5d,0x5e,0x41,0x5b,0x73,0x2b,
+    0x31,0x35,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x32,0x30,0x5d,0x3b,0x0a,0x44,0x5b,0x74,0x5d,0x3d,0x43,0x5b,0x62,0x5b,0x32,0x30,0x2b,0x73,0x5d,0x5d,0x5e,0x52,0x36,0x34,
+    0x28,0x43,0x5b,0x62,0x5b,0x35,0x2b,0x73,0x5d,0x5d,0x2c,0x31,0x2c,0x36,0x33,0x29,0x3b,0x0a,0x43,0x5b,0x74,0x5d,0x3d,0x52,0x36,0x34,0x28,0x41,0x5b,0x61,0x5b,0x74,
+    0x5d,0x5d,0x5e,0x44,0x5b,0x62,0x5b,0x74,0x5d,0x5d,0x2c,0x72,0x6f,0x5b,0x74,0x5d,0x5b,0x30,0x5d,0x2c,0x72,0x6f,0x5b,0x74,0x5d,0x5b,0x31,0x5d,0x29,0x3b,0x0a,0x41,
+    0x5b,0x64,0x5b,0x74,0x5d,0x5d,0x3d,0x43,0x5b,0x63,0x5b,0x74,0x5d,0x5b,0x30,0x5d,0x5d,0x5e,0x28,0x28,0x7e,0x43,0x5b,0x63,0x5b,0x74,0x5d,0x5b,0x31,0x5d,0x5d,0x29,
+    0x26,0x43,0x5b,0x63,0x5b,0x74,0x5d,0x5b,0x32,0x5d,0x5d,0x29,0x3b,0x0a,0x41,0x5b,0x74,0x5d,0x20,0x5e,0x3d,0x20,0x72,0x63,0x5b,0x28,0x74,0x3d,0x3d,0x30,0x29,0x3f,
+    0x30,0x3a,0x31,0x5d,0x5b,0x69,0x5d,0x3b,0x20,0x0a,0x7d,0x0a,0x69,0x66,0x28,0x74,0x3c,0x34,0x29,0x0a,0x7b,0x0a,0x68,0x61,0x73,0x68,0x65,0x73,0x2b,0x3d,0x67,0x2a,
+    0x28,0x33,0x32,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x29,0x3b,0x0a,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,0x74,0x5d,
+    0x3d,0x41,0x5b,0x74,0x5d,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x5f,0x5f,0x61,0x74,0x74,0x72,0x69,0x62,0x75,0x74,0x65,0x5f,0x5f,0x28,0x28,0x72,0x65,0x71,0x64,0x5f,0x77,
+    0x6f,0x72,0x6b,0x5f,0x67,0x72,0x6f,0x75,0x70,0x5f,0x73,0x69,0x7a,0x65,0x28,0x33,0x32,0x2c,0x31,0x2c,0x31,0x29,0x29,0x29,0x0a,0x5f,0x5f,0x6b,0x65,0x72,0x6e,0x65,
+    0x6c,0x20,0x76,0x6f,0x69,0x64,0x20,0x73,0x68,0x61,0x33,0x5f,0x69,0x6e,0x69,0x74,0x69,0x61,0x6c,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,
+    0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x20,0x69,0x6e,0x70,0x75,0x74,0x5f,0x64,0x61,0x74,0x61,0x2c,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,
+    0x69,0x6e,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x2c,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6e,0x6f,0x6e,0x63,0x65,0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,
+    0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x68,0x61,0x73,0x68,0x65,0x73,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,
+    0x74,0x33,0x32,0x5f,0x74,0x20,0x74,0x3d,0x67,0x65,0x74,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,
+    0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x3d,0x67,0x65,0x74,0x5f,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x69,0x66,0x28,0x74,0x3e,
+    0x3d,0x32,0x35,0x29,0x0a,0x72,0x65,0x74,0x75,0x72,0x6e,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x73,0x3d,0x74,0x20,
+    0x25,0x20,0x35,0x3b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x69,0x6e,0x70,0x75,0x74,0x3d,0x28,0x5f,
+    0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x29,0x28,0x69,0x6e,0x70,0x75,0x74,0x5f,0x64,0x61,0x74,0x61,0x29,0x3b,0x0a,
+    0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x41,0x5b,0x32,0x35,0x5d,0x3b,0x0a,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,
+    0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x43,0x5b,0x32,0x35,0x5d,0x3b,0x0a,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,
+    0x20,0x44,0x5b,0x32,0x35,0x5d,0x3b,0x0a,0x41,0x5b,0x74,0x5d,0x3d,0x28,0x74,0x3c,0x31,0x36,0x29,0x3f,0x69,0x6e,0x70,0x75,0x74,0x5b,0x74,0x5d,0x3a,0x30,0x3b,0x0a,
+    0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x6e,0x6f,0x6e,0x63,0x65,0x5f,0x70,0x6f,0x73,0x3d,0x28,0x5f,0x5f,0x6c,
+    0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x29,0x28,0x41,0x29,0x2b,0x39,0x3b,0x0a,0x6e,0x6f,0x6e,0x63,0x65,0x2b,0x3d,0x67,0x3b,0x0a,
+    0x6e,0x6f,0x6e,0x63,0x65,0x5f,0x70,0x6f,0x73,0x5b,0x30,0x5d,0x3d,0x28,0x6e,0x6f,0x6e,0x63,0x65,0x5f,0x70,0x6f,0x73,0x5b,0x30,0x5d,0x26,0x30,0x78,0x46,0x46,0x46,
+    0x46,0x46,0x46,0x55,0x29,0x7c,0x28,0x28,0x6e,0x6f,0x6e,0x63,0x65,0x26,0x30,0x78,0x46,0x46,0x29,0x3c,0x3c,0x32,0x34,0x29,0x3b,0x0a,0x6e,0x6f,0x6e,0x63,0x65,0x5f,
+    0x70,0x6f,0x73,0x5b,0x31,0x5d,0x3d,0x28,0x6e,0x6f,0x6e,0x63,0x65,0x5f,0x70,0x6f,0x73,0x5b,0x31,0x5d,0x26,0x30,0x78,0x46,0x46,0x30,0x30,0x30,0x30,0x30,0x30,0x55,
+    0x29,0x7c,0x28,0x6e,0x6f,0x6e,0x63,0x65,0x3e,0x3e,0x38,0x29,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x77,0x6f,0x72,0x64,0x49,0x6e,0x64,0x65,0x78,
+    0x3d,0x69,0x6e,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x3b,0x0a,0x63,0x6f,
+    0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x74,0x61,0x69,0x6c,0x5f,0x73,0x69,0x7a,0x65,0x3d,0x69,0x6e,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,
+    0x65,0x20,0x25,0x20,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x3b,0x0a,0x41,0x5b,0x77,0x6f,0x72,0x64,0x49,0x6e,0x64,0x65,
+    0x78,0x5d,0x20,0x5e,0x3d,0x20,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x28,0x28,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x28,0x30,0x78,0x30,
+    0x32,0x7c,0x28,0x31,0x3c,0x3c,0x32,0x29,0x29,0x29,0x3c,0x3c,0x28,0x74,0x61,0x69,0x6c,0x5f,0x73,0x69,0x7a,0x65,0x2a,0x38,0x29,0x29,0x3b,0x0a,0x41,0x5b,0x31,0x36,
+    0x5d,0x20,0x5e,0x3d,0x20,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x55,0x4c,0x3b,0x0a,0x23,0x70,0x72,0x61,0x67,
+    0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x28,0x52,0x4f,0x55,0x4e,0x44,0x53,0x29,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,
+    0x69,0x3c,0x52,0x4f,0x55,0x4e,0x44,0x53,0x3b,0x20,0x2b,0x2b,0x69,0x29,0x0a,0x7b,0x0a,0x43,0x5b,0x74,0x5d,0x3d,0x41,0x5b,0x73,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x35,
+    0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x31,0x30,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x31,0x35,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x32,0x30,0x5d,0x3b,0x0a,0x44,0x5b,0x74,0x5d,0x3d,
+    0x43,0x5b,0x62,0x5b,0x32,0x30,0x2b,0x73,0x5d,0x5d,0x5e,0x52,0x36,0x34,0x28,0x43,0x5b,0x62,0x5b,0x35,0x2b,0x73,0x5d,0x5d,0x2c,0x31,0x2c,0x36,0x33,0x29,0x3b,0x0a,
+    0x43,0x5b,0x74,0x5d,0x3d,0x52,0x36,0x34,0x28,0x41,0x5b,0x61,0x5b,0x74,0x5d,0x5d,0x5e,0x44,0x5b,0x62,0x5b,0x74,0x5d,0x5d,0x2c,0x72,0x6f,0x5b,0x74,0x5d,0x5b,0x30,
+    0x5d,0x2c,0x72,0x6f,0x5b,0x74,0x5d,0x5b,0x31,0x5d,0x29,0x3b,0x0a,0x41,0x5b,0x64,0x5b,0x74,0x5d,0x5d,0x3d,0x43,0x5b,0x63,0x5b,0x74,0x5d,0x5b,0x30,0x5d,0x5d,0x5e,
+    0x28,0x28,0x7e,0x43,0x5b,0x63,0x5b,0x74,0x5d,0x5b,0x31,0x5d,0x5d,0x29,0x26,0x43,0x5b,0x63,0x5b,0x74,0x5d,0x5b,0x32,0x5d,0x5d,0x29,0x3b,0x0a,0x41,0x5b,0x74,0x5d,
+    0x20,0x5e,0x3d,0x20,0x72,0x63,0x5b,0x28,0x74,0x3d,0x3d,0x30,0x29,0x3f,0x30,0x3a,0x31,0x5d,0x5b,0x69,0x5d,0x3b,0x20,0x0a,0x7d,0x0a,0x69,0x66,0x28,0x74,0x3c,0x34,
+    0x29,0x0a,0x7b,0x0a,0x68,0x61,0x73,0x68,0x65,0x73,0x2b,0x3d,0x67,0x2a,0x28,0x33,0x32,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,
+    0x74,0x29,0x29,0x3b,0x0a,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,0x74,0x5d,0x3d,0x41,0x5b,0x74,0x5d,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x00
+};
+
+} // namespace xmrig
diff --git a/src/backend/opencl/cl/astrobwt/salsa20.cl b/src/backend/opencl/cl/astrobwt/salsa20.cl
new file mode 100644
index 000000000..3c8431c74
--- /dev/null
+++ b/src/backend/opencl/cl/astrobwt/salsa20.cl
@@ -0,0 +1,153 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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 ROTATE(v,c) (rotate(v,(uint32_t)c))
+#define XOR(v,w) ((v) ^ (w))
+#define PLUS(v,w) ((v) + (w))
+
+__attribute__((reqd_work_group_size(SALSA20_GROUP_SIZE, 1, 1)))
+__kernel void Salsa20_XORKeyStream(__global const uint32_t* keys, __global uint32_t* outputs, __global uint32_t* output_sizes, uint32_t output_stride)
+{
+	const uint32_t t = get_local_id(0);
+	const uint32_t g = get_group_id(0);
+
+	// Put zeroes in the beginning
+	const uint64_t output_offset = g * ((uint64_t)output_stride) + 128;
+	{
+		__global uint32_t* p = outputs + (output_offset - 128) / sizeof(uint32_t);
+		for (uint32_t i = t; i < 128 / sizeof(uint32_t); i += SALSA20_GROUP_SIZE)
+			p[i] = 0;
+	}
+	
+	__global const uint32_t* k = keys + g * 8;
+	__global uint32_t* output = outputs + (output_offset + (t * 64)) / sizeof(uint32_t);
+	const uint32_t output_size = output_sizes[g];
+
+	const uint32_t j1 = k[0];
+	const uint32_t j2 = k[1];
+	const uint32_t j3 = k[2];
+	const uint32_t j4 = k[3];
+	const uint32_t j11 = k[4];
+	const uint32_t j12 = k[5];
+	const uint32_t j13 = k[6];
+	const uint32_t j14 = k[7];
+	const uint32_t j0 = 0x61707865U;
+	const uint32_t j5 = 0x3320646EU;
+	const uint32_t j10 = 0x79622D32U;
+	const uint32_t j15 = 0x6B206574U;
+	const uint32_t j6 = 0;
+	const uint32_t j7 = 0;
+	const uint32_t j8 = 0;
+	const uint32_t j9 = 0;
+
+	for (uint32_t i = t * 64; i < output_size; i += SALSA20_GROUP_SIZE * 64)
+	{
+		const uint32_t j8_1 = j8 + (i / 64);
+
+		uint32_t x0 = j0;
+		uint32_t x1 = j1;
+		uint32_t x2 = j2;
+		uint32_t x3 = j3;
+		uint32_t x4 = j4;
+		uint32_t x5 = j5;
+		uint32_t x6 = j6;
+		uint32_t x7 = j7;
+		uint32_t x8 = j8_1;
+		uint32_t x9 = j9;
+		uint32_t x10 = j10;
+		uint32_t x11 = j11;
+		uint32_t x12 = j12;
+		uint32_t x13 = j13;
+		uint32_t x14 = j14;
+		uint32_t x15 = j15;
+
+		#pragma unroll(5)
+		for (uint32_t j = 0; j < 10; ++j)
+		{
+			x4  = XOR( x4,ROTATE(PLUS( x0,x12), 7));
+			x8  = XOR( x8,ROTATE(PLUS( x4, x0), 9));
+			x12 = XOR(x12,ROTATE(PLUS( x8, x4),13));
+			x0  = XOR( x0,ROTATE(PLUS(x12, x8),18));
+			x9  = XOR( x9,ROTATE(PLUS( x5, x1), 7));
+			x13 = XOR(x13,ROTATE(PLUS( x9, x5), 9));
+			x1  = XOR( x1,ROTATE(PLUS(x13, x9),13));
+			x5  = XOR( x5,ROTATE(PLUS( x1,x13),18));
+			x14 = XOR(x14,ROTATE(PLUS(x10, x6), 7));
+			x2  = XOR( x2,ROTATE(PLUS(x14,x10), 9));
+			x6  = XOR( x6,ROTATE(PLUS( x2,x14),13));
+			x10 = XOR(x10,ROTATE(PLUS( x6, x2),18));
+			x3  = XOR( x3,ROTATE(PLUS(x15,x11), 7));
+			x7  = XOR( x7,ROTATE(PLUS( x3,x15), 9));
+			x11 = XOR(x11,ROTATE(PLUS( x7, x3),13));
+			x15 = XOR(x15,ROTATE(PLUS(x11, x7),18));
+			x1  = XOR( x1,ROTATE(PLUS( x0, x3), 7));
+			x2  = XOR( x2,ROTATE(PLUS( x1, x0), 9));
+			x3  = XOR( x3,ROTATE(PLUS( x2, x1),13));
+			x0  = XOR( x0,ROTATE(PLUS( x3, x2),18));
+			x6  = XOR( x6,ROTATE(PLUS( x5, x4), 7));
+			x7  = XOR( x7,ROTATE(PLUS( x6, x5), 9));
+			x4  = XOR( x4,ROTATE(PLUS( x7, x6),13));
+			x5  = XOR( x5,ROTATE(PLUS( x4, x7),18));
+			x11 = XOR(x11,ROTATE(PLUS(x10, x9), 7));
+			x8  = XOR( x8,ROTATE(PLUS(x11,x10), 9));
+			x9  = XOR( x9,ROTATE(PLUS( x8,x11),13));
+			x10 = XOR(x10,ROTATE(PLUS( x9, x8),18));
+			x12 = XOR(x12,ROTATE(PLUS(x15,x14), 7));
+			x13 = XOR(x13,ROTATE(PLUS(x12,x15), 9));
+			x14 = XOR(x14,ROTATE(PLUS(x13,x12),13));
+			x15 = XOR(x15,ROTATE(PLUS(x14,x13),18));
+		}
+
+		output[0]  = PLUS(x0, j0);
+		output[1]  = PLUS(x1, j1);
+		output[2]  = PLUS(x2, j2);
+		output[3]  = PLUS(x3, j3);
+		output[4]  = PLUS(x4, j4);
+		output[5]  = PLUS(x5, j5);
+		output[6]  = PLUS(x6, j6);
+		output[7]  = PLUS(x7, j7);
+		output[8]  = PLUS(x8, j8_1);
+		output[9]  = PLUS(x9, j9);
+		output[10] = PLUS(x10,j10);
+		output[11] = PLUS(x11,j11);
+		output[12] = PLUS(x12,j12);
+		output[13] = PLUS(x13,j13);
+		output[14] = PLUS(x14,j14);
+		output[15] = PLUS(x15,j15);
+
+		output += (SALSA20_GROUP_SIZE * 64) / sizeof(uint32_t);
+	}
+
+	barrier(CLK_GLOBAL_MEM_FENCE);
+
+	// Put zeroes after output's end
+	if (t < 16)
+	{
+		__global uint32_t* p = outputs + (output_offset + output_size + 3) / sizeof(uint32_t);
+		p[t] = 0;
+	}
+
+	if ((t == 0) && (output_size & 3))
+		outputs[(output_offset + output_size) / sizeof(uint32_t)] &= 0xFFFFFFFFU >> ((4 - (output_size & 3)) << 3);
+}
diff --git a/src/backend/opencl/cl/astrobwt/sha3.cl b/src/backend/opencl/cl/astrobwt/sha3.cl
new file mode 100644
index 000000000..8c2ee24de
--- /dev/null
+++ b/src/backend/opencl/cl/astrobwt/sha3.cl
@@ -0,0 +1,198 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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 ROUNDS 24 
+#define R64(a,b,c) (((a) << b) | ((a) >> c))
+
+__constant const uint64_t rc[2][ROUNDS] = {
+	{0x0000000000000001UL, 0x0000000000008082UL, 0x800000000000808AUL,
+	0x8000000080008000UL, 0x000000000000808BUL, 0x0000000080000001UL,
+	0x8000000080008081UL, 0x8000000000008009UL, 0x000000000000008AUL,
+	0x0000000000000088UL, 0x0000000080008009UL, 0x000000008000000AUL,
+	0x000000008000808BUL, 0x800000000000008BUL, 0x8000000000008089UL,
+	0x8000000000008003UL, 0x8000000000008002UL, 0x8000000000000080UL,
+	0x000000000000800AUL, 0x800000008000000AUL, 0x8000000080008081UL,
+	0x8000000000008080UL, 0x0000000080000001UL, 0x8000000080008008UL},
+	{0UL, 0UL, 0UL, 0UL, 0UL, 0UL, 0UL, 0UL,
+	0UL, 0UL, 0UL, 0UL, 0UL, 0UL, 0UL, 0UL,
+	0UL, 0UL, 0UL, 0UL, 0UL, 0UL, 0UL, 0UL}
+};
+
+
+__constant const int ro[25][2] = {
+	{ 0,64}, {44,20}, {43,21}, {21,43}, {14,50},
+	{ 1,63}, { 6,58}, {25,39}, { 8,56}, {18,46},
+	{62, 2}, {55, 9}, {39,25}, {41,23}, { 2,62},
+	{28,36}, {20,44}, { 3,61}, {45,19}, {61, 3},
+	{27,37}, {36,28}, {10,54}, {15,49}, {56, 8}
+};
+
+__constant const int a[25] = {
+	0,  6, 12, 18, 24,
+	1,  7, 13, 19, 20,
+	2,  8, 14, 15, 21,
+	3,  9, 10, 16, 22,
+	4,  5, 11, 17, 23
+};
+
+__constant const int b[25] = {
+	0,  1,  2,  3, 4,
+	1,  2,  3,  4, 0,
+	2,  3,  4,  0, 1,
+	3,  4,  0,  1, 2,
+	4,  0,  1,  2, 3
+};
+
+__constant const int c[25][3] = {
+	{ 0, 1, 2}, { 1, 2, 3}, { 2, 3, 4}, { 3, 4, 0}, { 4, 0, 1},
+	{ 5, 6, 7}, { 6, 7, 8}, { 7, 8, 9}, { 8, 9, 5}, { 9, 5, 6},
+	{10,11,12}, {11,12,13}, {12,13,14}, {13,14,10}, {14,10,11},
+	{15,16,17}, {16,17,18}, {17,18,19}, {18,19,15}, {19,15,16},
+	{20,21,22}, {21,22,23}, {22,23,24}, {23,24,20}, {24,20,21}
+};
+
+__constant const int d[25] = {
+	0,  1,  2,  3,  4,
+	10, 11, 12, 13, 14,
+	20, 21, 22, 23, 24,
+	5,  6,  7,  8,  9,
+	15, 16, 17, 18, 19
+};
+
+__attribute__((reqd_work_group_size(32, 1, 1)))
+__kernel void sha3(__global const uint8_t* inputs, __global const uint32_t* input_sizes, uint32_t input_stride, __global uint64_t* hashes)
+{
+	const uint32_t t = get_local_id(0);
+	const uint32_t g = get_group_id(0);
+
+	if (t >= 25)
+		return;
+
+	const uint32_t s = t % 5;
+
+	const uint64_t input_offset = ((uint64_t)input_stride) * g;
+	__global uint64_t* input = (__global uint64_t*)(inputs + input_offset);
+	const uint32_t input_size = input_sizes[g] + 1;
+
+	__local uint64_t A[25];
+	__local uint64_t C[25];
+	__local uint64_t D[25];
+
+	A[t] = 0;
+
+	const uint32_t words = input_size / sizeof(uint64_t);
+	const uint32_t tail_size = input_size % sizeof(uint64_t);
+
+	uint32_t wordIndex = 0;
+	for (uint32_t i = 0; i < words; ++i, ++input)
+	{
+		A[wordIndex] ^= *input;
+		++wordIndex;
+		if (wordIndex == 17)
+		{
+			#pragma unroll(ROUNDS)
+			for (int i = 0; i < ROUNDS; ++i)
+			{
+				C[t] = A[s] ^ A[s+5] ^ A[s+10] ^ A[s+15] ^ A[s+20];
+				D[t] = C[b[20+s]] ^ R64(C[b[5+s]], 1, 63);
+				C[t] = R64(A[a[t]] ^ D[b[t]], ro[t][0], ro[t][1]);
+				A[d[t]] = C[c[t][0]] ^ ((~C[c[t][1]]) & C[c[t][2]]);
+				A[t] ^= rc[(t == 0) ? 0 : 1][i]; 
+			}
+			wordIndex = 0;
+		}
+	}
+
+	uint64_t tail = 0;
+	__global const uint8_t* p = (__global const uint8_t*)input;
+	for (uint32_t i = 0; i < tail_size; ++i)
+	{
+		tail |= (uint64_t)(p[i]) << (i * 8);
+	}
+	A[wordIndex] ^= tail ^ ((uint64_t)(((uint64_t)(0x02 | (1 << 2))) << (tail_size * 8)));
+	A[16] ^= 0x8000000000000000UL;
+
+	#pragma unroll(1)
+	for (int i = 0; i < ROUNDS; ++i)
+	{
+		C[t] = A[s] ^ A[s+5] ^ A[s+10] ^ A[s+15] ^ A[s+20];
+		D[t] = C[b[20+s]] ^ R64(C[b[5+s]], 1, 63);
+		C[t] = R64(A[a[t]] ^ D[b[t]], ro[t][0], ro[t][1]);
+		A[d[t]] = C[c[t][0]] ^ ((~C[c[t][1]]) & C[c[t][2]]);
+		A[t] ^= rc[(t == 0) ? 0 : 1][i]; 
+	}
+
+	if (t < 4)
+	{
+		hashes += g * (32 / sizeof(uint64_t));
+		hashes[t] = A[t];
+	}
+}
+
+__attribute__((reqd_work_group_size(32, 1, 1)))
+__kernel void sha3_initial(__global const uint8_t* input_data, uint32_t input_size, uint32_t nonce, __global uint64_t* hashes)
+{
+	const uint32_t t = get_local_id(0);
+	const uint32_t g = get_group_id(0);
+
+	if (t >= 25)
+		return;
+
+	const uint32_t s = t % 5;
+
+	__global uint64_t* input = (__global uint64_t*)(input_data);
+
+	__local uint64_t A[25];
+	__local uint64_t C[25];
+	__local uint64_t D[25];
+
+	A[t] = (t < 16) ? input[t] : 0;
+
+	__local uint32_t* nonce_pos = (__local uint32_t*)(A) + 9;
+	nonce += g;
+	nonce_pos[0] = (nonce_pos[0] & 0xFFFFFFU) | ((nonce & 0xFF) << 24);
+	nonce_pos[1] = (nonce_pos[1] & 0xFF000000U) | (nonce >> 8);
+
+	uint32_t wordIndex = input_size / sizeof(uint64_t);
+	const uint32_t tail_size = input_size % sizeof(uint64_t);
+
+	A[wordIndex] ^= (uint64_t)(((uint64_t)(0x02 | (1 << 2))) << (tail_size * 8));
+	A[16] ^= 0x8000000000000000UL;
+
+	#pragma unroll(ROUNDS)
+	for (int i = 0; i < ROUNDS; ++i)
+	{
+		C[t] = A[s] ^ A[s+5] ^ A[s+10] ^ A[s+15] ^ A[s+20];
+		D[t] = C[b[20+s]] ^ R64(C[b[5+s]], 1, 63);
+		C[t] = R64(A[a[t]] ^ D[b[t]], ro[t][0], ro[t][1]);
+		A[d[t]] = C[c[t][0]] ^ ((~C[c[t][1]]) & C[c[t][2]]);
+		A[t] ^= rc[(t == 0) ? 0 : 1][i]; 
+	}
+
+	if (t < 4)
+	{
+		hashes += g * (32 / sizeof(uint64_t));
+		hashes[t] = A[t];
+	}
+}
diff --git a/src/backend/opencl/generators/ocl_generic_astrobwt_generator.cpp b/src/backend/opencl/generators/ocl_generic_astrobwt_generator.cpp
new file mode 100644
index 000000000..95dffd5b8
--- /dev/null
+++ b/src/backend/opencl/generators/ocl_generic_astrobwt_generator.cpp
@@ -0,0 +1,64 @@
+/* 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      Lee Clagett <https://github.com/vtnerd>
+ * Copyright 2018-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 XMRig       <https://github.com/xmrig>, <support@xmrig.com>
+ *
+ *   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/OclThreads.h"
+#include "backend/opencl/wrappers/OclDevice.h"
+#include "base/crypto/Algorithm.h"
+#include "crypto/randomx/randomx.h"
+#include "crypto/rx/RxAlgo.h"
+
+
+namespace xmrig {
+
+
+bool ocl_generic_astrobwt_generator(const OclDevice &device, const Algorithm &algorithm, OclThreads &threads)
+{
+    if (algorithm.family() != Algorithm::ASTROBWT) {
+        return false;
+    }
+
+    const size_t mem = device.globalMemSize();
+
+    uint32_t per_thread_mem = 10 << 20;
+    uint32_t intensity = static_cast<uint32_t>((mem - (128 << 20)) / per_thread_mem / 2);
+
+    intensity &= ~63U;
+
+    if (!intensity) {
+        return false;
+    }
+
+    if (intensity > 256) {
+        intensity = 256;
+    }
+
+    threads.add(OclThread(device.index(), intensity, 2));
+
+    return true;
+}
+
+
+} // namespace xmrig
diff --git a/src/backend/opencl/interfaces/IOclRunner.h b/src/backend/opencl/interfaces/IOclRunner.h
index 0b47bcd29..74f02c6cb 100644
--- a/src/backend/opencl/interfaces/IOclRunner.h
+++ b/src/backend/opencl/interfaces/IOclRunner.h
@@ -59,6 +59,8 @@ public:
     virtual const OclLaunchData &data() const               = 0;
     virtual size_t intensity() const                        = 0;
     virtual size_t threadId() const                         = 0;
+    virtual uint32_t roundSize() const                      = 0;
+    virtual uint32_t processedHashes() const                = 0;
     virtual uint32_t deviceIndex() const                    = 0;
     virtual void build()                                    = 0;
     virtual void init()                                     = 0;
diff --git a/src/backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.cpp b/src/backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.cpp
new file mode 100644
index 000000000..95f65d327
--- /dev/null
+++ b/src/backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.cpp
@@ -0,0 +1,42 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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/astrobwt/AstroBWT_FilterKernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+void xmrig::AstroBWT_FilterKernel::enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size)
+{
+    enqueueNDRange(queue, 1, nullptr, &threads, &workgroup_size);
+}
+
+
+void xmrig::AstroBWT_FilterKernel::setArgs(uint32_t nonce, uint32_t bwt_max_size, cl_mem hashes, cl_mem filtered_hashes)
+{
+    setArg(0, sizeof(uint32_t), &nonce);
+    setArg(1, sizeof(uint32_t), &bwt_max_size);
+    setArg(2, sizeof(cl_mem), &hashes);
+    setArg(3, sizeof(cl_mem), &filtered_hashes);
+}
diff --git a/src/backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.h b/src/backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.h
new file mode 100644
index 000000000..60a748de0
--- /dev/null
+++ b/src/backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.h
@@ -0,0 +1,48 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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_ASTROBWT_FILTERKERNEL_H
+#define XMRIG_ASTROBWT_FILTERKERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+class AstroBWT_FilterKernel : public OclKernel
+{
+public:
+    inline AstroBWT_FilterKernel(cl_program program) : OclKernel(program, "filter") {}
+
+    void enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size);
+    void setArgs(uint32_t nonce, uint32_t bwt_max_size, cl_mem hashes, cl_mem filtered_hashes);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_ASTROBWT_FILTERKERNEL_H */
diff --git a/src/backend/opencl/kernels/astrobwt/AstroBWT_FindSharesKernel.cpp b/src/backend/opencl/kernels/astrobwt/AstroBWT_FindSharesKernel.cpp
new file mode 100644
index 000000000..4f3d8aee1
--- /dev/null
+++ b/src/backend/opencl/kernels/astrobwt/AstroBWT_FindSharesKernel.cpp
@@ -0,0 +1,47 @@
+/* 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/astrobwt/AstroBWT_FindSharesKernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+void xmrig::AstroBWT_FindSharesKernel::enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size)
+{
+    enqueueNDRange(queue, 1, nullptr, &threads, &workgroup_size);
+}
+
+
+void xmrig::AstroBWT_FindSharesKernel::setArgs(cl_mem hashes, cl_mem filtered_hashes, cl_mem shares)
+{
+    setArg(0, sizeof(cl_mem), &hashes);
+    setArg(1, sizeof(cl_mem), &filtered_hashes);
+    setArg(3, sizeof(cl_mem), &shares);
+}
+
+
+void xmrig::AstroBWT_FindSharesKernel::setTarget(uint64_t target)
+{
+    setArg(2, sizeof(uint64_t), &target);
+}
diff --git a/src/backend/opencl/kernels/astrobwt/AstroBWT_FindSharesKernel.h b/src/backend/opencl/kernels/astrobwt/AstroBWT_FindSharesKernel.h
new file mode 100644
index 000000000..6645cc56b
--- /dev/null
+++ b/src/backend/opencl/kernels/astrobwt/AstroBWT_FindSharesKernel.h
@@ -0,0 +1,49 @@
+/* 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_ASTROBWTFINDSHARESKERNEL_H
+#define XMRIG_ASTROBWTFINDSHARESKERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+class AstroBWT_FindSharesKernel : public OclKernel
+{
+public:
+    inline AstroBWT_FindSharesKernel(cl_program program) : OclKernel(program, "find_shares") {}
+
+    void enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size);
+    void setArgs(cl_mem hashes, cl_mem filtered_hashes, cl_mem shares);
+    void setTarget(uint64_t target);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_ASTROBWTFINDSHARESKERNEL_H */
diff --git a/src/backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.cpp b/src/backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.cpp
new file mode 100644
index 000000000..c08562a6f
--- /dev/null
+++ b/src/backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.cpp
@@ -0,0 +1,44 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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/astrobwt/AstroBWT_MainKernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+void xmrig::AstroBWT_MainKernel::enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size)
+{
+    const size_t gthreads       = threads * workgroup_size;
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &workgroup_size);
+}
+
+
+void xmrig::AstroBWT_MainKernel::setArgs(cl_mem inputs, cl_mem input_sizes, uint32_t input_stride, cl_mem indices, cl_mem tmp_indices)
+{
+    setArg(0, sizeof(cl_mem), &inputs);
+    setArg(1, sizeof(cl_mem), &input_sizes);
+    setArg(2, sizeof(uint32_t), &input_stride);
+    setArg(3, sizeof(cl_mem), &indices);
+    setArg(4, sizeof(cl_mem), &tmp_indices);
+}
diff --git a/src/backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.h b/src/backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.h
new file mode 100644
index 000000000..812256ebc
--- /dev/null
+++ b/src/backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.h
@@ -0,0 +1,48 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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_ASTROBWT_MAINKERNEL_H
+#define XMRIG_ASTROBWT_MAINKERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+class AstroBWT_MainKernel : public OclKernel
+{
+public:
+    inline AstroBWT_MainKernel(cl_program program) : OclKernel(program, "BWT") {}
+
+    void enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size);
+    void setArgs(cl_mem inputs, cl_mem input_sizes, uint32_t input_stride, cl_mem indices, cl_mem tmp_indices);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_ASTROBWT_MAINKERNEL_H */
diff --git a/src/backend/opencl/kernels/astrobwt/AstroBWT_PrepareBatch2Kernel.cpp b/src/backend/opencl/kernels/astrobwt/AstroBWT_PrepareBatch2Kernel.cpp
new file mode 100644
index 000000000..962e48218
--- /dev/null
+++ b/src/backend/opencl/kernels/astrobwt/AstroBWT_PrepareBatch2Kernel.cpp
@@ -0,0 +1,41 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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/astrobwt/AstroBWT_PrepareBatch2Kernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+void xmrig::AstroBWT_PrepareBatch2Kernel::enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size)
+{
+    enqueueNDRange(queue, 1, nullptr, &threads, &workgroup_size);
+}
+
+
+void xmrig::AstroBWT_PrepareBatch2Kernel::setArgs(cl_mem hashes, cl_mem filtered_hashes, cl_mem bwt_data_sizes)
+{
+    setArg(0, sizeof(cl_mem), &hashes);
+    setArg(1, sizeof(cl_mem), &filtered_hashes);
+    setArg(2, sizeof(cl_mem), &bwt_data_sizes);
+}
diff --git a/src/backend/opencl/kernels/astrobwt/AstroBWT_PrepareBatch2Kernel.h b/src/backend/opencl/kernels/astrobwt/AstroBWT_PrepareBatch2Kernel.h
new file mode 100644
index 000000000..8dd1f9836
--- /dev/null
+++ b/src/backend/opencl/kernels/astrobwt/AstroBWT_PrepareBatch2Kernel.h
@@ -0,0 +1,48 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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_ASTROBWT_PREPAREBATCH2KERNEL_H
+#define XMRIG_ASTROBWT_PREPAREBATCH2KERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+class AstroBWT_PrepareBatch2Kernel : public OclKernel
+{
+public:
+    inline AstroBWT_PrepareBatch2Kernel(cl_program program) : OclKernel(program, "prepare_batch2") {}
+
+    void enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size);
+    void setArgs(cl_mem hashes, cl_mem filtered_hashes, cl_mem bwt_data_sizes);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_ASTROBWT_PREPAREBATCH2KERNEL_H */
diff --git a/src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3InitialKernel.cpp b/src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3InitialKernel.cpp
new file mode 100644
index 000000000..ac80ae101
--- /dev/null
+++ b/src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3InitialKernel.cpp
@@ -0,0 +1,44 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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/astrobwt/AstroBWT_SHA3InitialKernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+void xmrig::AstroBWT_SHA3InitialKernel::enqueue(cl_command_queue queue, size_t threads)
+{
+    const size_t workgroup_size = 32;
+    const size_t gthreads       = threads * workgroup_size;
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &workgroup_size);
+}
+
+
+void xmrig::AstroBWT_SHA3InitialKernel::setArgs(cl_mem input, uint32_t input_size, uint32_t nonce, cl_mem output_salsa20_keys)
+{
+    setArg(0, sizeof(cl_mem), &input);
+    setArg(1, sizeof(uint32_t), &input_size);
+    setArg(2, sizeof(uint32_t), &nonce);
+    setArg(3, sizeof(cl_mem), &output_salsa20_keys);
+}
diff --git a/src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3InitialKernel.h b/src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3InitialKernel.h
new file mode 100644
index 000000000..ad6486c58
--- /dev/null
+++ b/src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3InitialKernel.h
@@ -0,0 +1,48 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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_ASTROBWT_SHA3INITIALKERNEL_H
+#define XMRIG_ASTROBWT_SHA3INITIALKERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+class AstroBWT_SHA3InitialKernel : public OclKernel
+{
+public:
+    inline AstroBWT_SHA3InitialKernel(cl_program program) : OclKernel(program, "sha3_initial") {}
+
+    void enqueue(cl_command_queue queue, size_t threads);
+    void setArgs(cl_mem input, uint32_t input_size, uint32_t nonce, cl_mem output_salsa20_keys);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_ASTROBWT_SHA3INITIALKERNEL_H */
diff --git a/src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.cpp b/src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.cpp
new file mode 100644
index 000000000..323505cb1
--- /dev/null
+++ b/src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.cpp
@@ -0,0 +1,44 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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/astrobwt/AstroBWT_SHA3Kernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+void xmrig::AstroBWT_SHA3Kernel::enqueue(cl_command_queue queue, size_t threads)
+{
+    const size_t workgroup_size = 32;
+    const size_t gthreads       = threads * workgroup_size;
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &workgroup_size);
+}
+
+
+void xmrig::AstroBWT_SHA3Kernel::setArgs(cl_mem inputs, cl_mem input_sizes, uint32_t input_stride, cl_mem output_salsa20_keys)
+{
+    setArg(0, sizeof(cl_mem), &inputs);
+    setArg(1, sizeof(cl_mem), &input_sizes);
+    setArg(2, sizeof(uint32_t), &input_stride);
+    setArg(3, sizeof(cl_mem), &output_salsa20_keys);
+}
diff --git a/src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.h b/src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.h
new file mode 100644
index 000000000..62dc1a1a4
--- /dev/null
+++ b/src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.h
@@ -0,0 +1,48 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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_ASTROBWT_SHA3KERNEL_H
+#define XMRIG_ASTROBWT_SHA3KERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+class AstroBWT_SHA3Kernel : public OclKernel
+{
+public:
+    inline AstroBWT_SHA3Kernel(cl_program program) : OclKernel(program, "sha3") {}
+
+    void enqueue(cl_command_queue queue, size_t threads);
+    void setArgs(cl_mem inputs, cl_mem input_sizes, uint32_t input_stride, cl_mem output_salsa20_keys);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_ASTROBWT_SHA3KERNEL_H */
diff --git a/src/backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.cpp b/src/backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.cpp
new file mode 100644
index 000000000..8a0ae9ae5
--- /dev/null
+++ b/src/backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.cpp
@@ -0,0 +1,43 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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/astrobwt/AstroBWT_Salsa20Kernel.h"
+#include "backend/opencl/wrappers/OclLib.h"
+
+
+void xmrig::AstroBWT_Salsa20Kernel::enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size)
+{
+    const size_t gthreads       = threads * workgroup_size;
+    enqueueNDRange(queue, 1, nullptr, &gthreads, &workgroup_size);
+}
+
+
+void xmrig::AstroBWT_Salsa20Kernel::setArgs(cl_mem salsa20_keys, cl_mem outputs, cl_mem output_sizes, uint32_t output_stride)
+{
+    setArg(0, sizeof(cl_mem), &salsa20_keys);
+    setArg(1, sizeof(cl_mem), &outputs);
+    setArg(2, sizeof(cl_mem), &output_sizes);
+    setArg(3, sizeof(uint32_t), &output_stride);
+}
diff --git a/src/backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.h b/src/backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.h
new file mode 100644
index 000000000..945b1fd6c
--- /dev/null
+++ b/src/backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.h
@@ -0,0 +1,48 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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_ASTROBWT_SALSA20KERNEL_H
+#define XMRIG_ASTROBWT_SALSA20KERNEL_H
+
+
+#include "backend/opencl/wrappers/OclKernel.h"
+
+
+namespace xmrig {
+
+
+class AstroBWT_Salsa20Kernel : public OclKernel
+{
+public:
+    inline AstroBWT_Salsa20Kernel(cl_program program) : OclKernel(program, "Salsa20_XORKeyStream") {}
+
+    void enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size);
+    void setArgs(cl_mem salsa20_keys, cl_mem outputs, cl_mem output_sizes, uint32_t output_stride);
+};
+
+
+} // namespace xmrig
+
+
+#endif /* XMRIG_ASTROBWT_SALSA20KERNEL_H */
diff --git a/src/backend/opencl/opencl.cmake b/src/backend/opencl/opencl.cmake
index 535e1a731..ebee46734 100644
--- a/src/backend/opencl/opencl.cmake
+++ b/src/backend/opencl/opencl.cmake
@@ -100,6 +100,31 @@ if (WITH_OPENCL)
              )
     endif()
 
+    if (WITH_ASTROBWT)
+        list(APPEND HEADERS_BACKEND_OPENCL
+             src/backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.h
+             src/backend/opencl/kernels/astrobwt/AstroBWT_FindSharesKernel.h
+             src/backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.h
+             src/backend/opencl/kernels/astrobwt/AstroBWT_PrepareBatch2Kernel.h
+             src/backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.h
+             src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3InitialKernel.h
+             src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.h
+             src/backend/opencl/runners/OclAstroBWTRunner.h
+             )
+
+        list(APPEND SOURCES_BACKEND_OPENCL
+             src/backend/opencl/generators/ocl_generic_astrobwt_generator.cpp
+             src/backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.cpp
+             src/backend/opencl/kernels/astrobwt/AstroBWT_FindSharesKernel.cpp
+             src/backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.cpp
+             src/backend/opencl/kernels/astrobwt/AstroBWT_PrepareBatch2Kernel.cpp
+             src/backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.cpp
+             src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3InitialKernel.cpp
+             src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.cpp
+             src/backend/opencl/runners/OclAstroBWTRunner.cpp
+             )
+    endif()
+
     if (WITH_CN_GPU AND CMAKE_SIZEOF_VOID_P EQUAL 8)
         list(APPEND HEADERS_BACKEND_OPENCL
              src/backend/opencl/kernels/Cn00RyoKernel.h
diff --git a/src/backend/opencl/runners/OclAstroBWTRunner.cpp b/src/backend/opencl/runners/OclAstroBWTRunner.cpp
new file mode 100644
index 000000000..94dee4257
--- /dev/null
+++ b/src/backend/opencl/runners/OclAstroBWTRunner.cpp
@@ -0,0 +1,211 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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/runners/OclAstroBWTRunner.h"
+#include "backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.h"
+#include "backend/opencl/kernels/astrobwt/AstroBWT_FindSharesKernel.h"
+#include "backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.h"
+#include "backend/opencl/kernels/astrobwt/AstroBWT_PrepareBatch2Kernel.h"
+#include "backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.h"
+#include "backend/opencl/kernels/astrobwt/AstroBWT_SHA3InitialKernel.h"
+#include "backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.h"
+#include "backend/opencl/OclLaunchData.h"
+#include "backend/opencl/wrappers/OclLib.h"
+#include "base/io/log/Log.h"
+#include "base/net/stratum/Job.h"
+
+
+constexpr int STAGE1_SIZE = 147253;
+constexpr uint32_t STAGE1_DATA_STRIDE = (STAGE1_SIZE + 256 + 255) & ~255U;
+
+
+xmrig::OclAstroBWTRunner::OclAstroBWTRunner(size_t index, const OclLaunchData &data) : OclBaseRunner(index, data)
+{
+    switch (data.device.type())
+    {
+    case OclDevice::Baffin:
+    case OclDevice::Polaris:
+    case OclDevice::Lexa:
+    case OclDevice::Vega_10:
+    case OclDevice::Vega_20:
+    case OclDevice::Raven:
+        m_workgroup_size = 64;
+        break;
+
+    default:
+        m_workgroup_size = 32;
+        break;
+    }
+
+    m_options += " -DSALSA20_GROUP_SIZE=" + std::to_string(m_workgroup_size);
+    m_options += " -DBWT_GROUP_SIZE="     + std::to_string(m_workgroup_size);
+
+    m_bwt_allocation_size = static_cast<uint64_t>(m_intensity) * BWT_DATA_STRIDE;
+    m_batch_size1 = static_cast<uint32_t>(m_bwt_allocation_size / STAGE1_DATA_STRIDE) & ~255U;
+
+    m_bwt_data_sizes_host = new uint32_t[m_batch_size1];
+}
+
+
+xmrig::OclAstroBWTRunner::~OclAstroBWTRunner()
+{
+    delete m_sha3_initial_kernel;
+    delete m_sha3_kernel;
+    delete m_salsa20_kernel;
+    delete m_bwt_kernel;
+    delete m_filter_kernel;
+    delete m_prepare_batch2_kernel;
+    delete m_find_shares_kernel;
+
+    OclLib::release(m_salsa20_keys);
+    OclLib::release(m_bwt_data);
+    OclLib::release(m_bwt_data_sizes);
+    OclLib::release(m_indices);
+    OclLib::release(m_tmp_indices);
+    OclLib::release(m_filtered_hashes);
+
+    delete m_bwt_data_sizes_host;
+}
+
+
+size_t xmrig::OclAstroBWTRunner::bufferSize() const
+{
+    return OclBaseRunner::bufferSize() +
+        align(m_batch_size1 * 32) +                    // m_salsa20_keys
+        align(m_bwt_allocation_size) +                 // m_bwt_data
+        align(m_batch_size1 * 4) +                     // m_bwt_data_sizes
+        align(m_bwt_allocation_size * 8) +             // m_indices
+        align(m_bwt_allocation_size * 8) +             // m_tmp_indices
+        align((m_batch_size1 + m_intensity) * 36 + 4); // m_filtered_hashes
+}
+
+
+void xmrig::OclAstroBWTRunner::run(uint32_t nonce, uint32_t *hashOutput)
+{
+    m_sha3_initial_kernel->setArg(2, sizeof(nonce), &nonce);
+    m_salsa20_kernel->setArg(3, sizeof(STAGE1_DATA_STRIDE), &STAGE1_DATA_STRIDE);
+    m_bwt_kernel->setArg(2, sizeof(STAGE1_DATA_STRIDE), &STAGE1_DATA_STRIDE);
+
+    const uint32_t t = STAGE1_DATA_STRIDE * 8;
+    m_sha3_kernel->setArg(2, sizeof(t), &t);
+    m_filter_kernel->setArg(0, sizeof(nonce), &nonce);
+
+    const uint32_t zero = 0;
+    enqueueWriteBuffer(m_output, CL_FALSE, sizeof(cl_uint) * 0xFF, sizeof(uint32_t), &zero);
+
+    m_sha3_initial_kernel->enqueue(m_queue, m_batch_size1);
+
+    for (uint32_t i = 0; i < m_batch_size1; ++i)
+        m_bwt_data_sizes_host[i] = STAGE1_SIZE;
+
+    enqueueWriteBuffer(m_bwt_data_sizes, CL_FALSE, 0, m_batch_size1 * sizeof(uint32_t), m_bwt_data_sizes_host);
+
+    m_salsa20_kernel->enqueue(m_queue, m_batch_size1, m_workgroup_size);
+    m_bwt_kernel->enqueue(m_queue, m_batch_size1, m_workgroup_size);
+    m_sha3_kernel->enqueue(m_queue, m_batch_size1);
+    m_filter_kernel->enqueue(m_queue, m_batch_size1, m_workgroup_size);
+
+    uint32_t num_filtered_hashes = 0;
+    enqueueReadBuffer(m_filtered_hashes, CL_TRUE, 0, sizeof(num_filtered_hashes), &num_filtered_hashes);
+
+    m_processedHashes = 0;
+    while (num_filtered_hashes >= m_intensity)
+    {
+        num_filtered_hashes -= m_intensity;
+        m_processedHashes += m_intensity;
+
+        m_salsa20_kernel->setArg(3, sizeof(BWT_DATA_STRIDE), &BWT_DATA_STRIDE);
+        m_bwt_kernel->setArg(2, sizeof(BWT_DATA_STRIDE), &BWT_DATA_STRIDE);
+
+        const uint32_t t = BWT_DATA_STRIDE * 8;
+        m_sha3_kernel->setArg(2, sizeof(t), &t);
+
+        m_prepare_batch2_kernel->enqueue(m_queue, m_intensity, m_workgroup_size);
+        m_salsa20_kernel->enqueue(m_queue, m_intensity, m_workgroup_size);
+        m_bwt_kernel->enqueue(m_queue, m_intensity, m_workgroup_size);
+        m_sha3_kernel->enqueue(m_queue, m_intensity);
+
+        m_find_shares_kernel->enqueue(m_queue, m_intensity, m_workgroup_size);
+
+        finalize(hashOutput);
+
+        OclLib::finish(m_queue);
+    }
+}
+
+
+void xmrig::OclAstroBWTRunner::set(const Job &job, uint8_t *blob)
+{
+    if (job.size() > (Job::kMaxBlobSize - 4)) {
+        throw std::length_error("job size too big");
+    }
+
+    if (job.size() < Job::kMaxBlobSize) {
+        memset(blob + job.size(), 0, Job::kMaxBlobSize - job.size());
+    }
+
+    enqueueWriteBuffer(m_input, CL_TRUE, 0, Job::kMaxBlobSize, blob);
+
+    m_sha3_initial_kernel->setArgs(m_input, static_cast<uint32_t>(job.size()), *job.nonce(), m_salsa20_keys);
+    m_salsa20_kernel->setArgs(m_salsa20_keys, m_bwt_data, m_bwt_data_sizes, STAGE1_DATA_STRIDE);
+    m_bwt_kernel->setArgs(m_bwt_data, m_bwt_data_sizes, STAGE1_DATA_STRIDE, m_indices, m_tmp_indices);
+    m_sha3_kernel->setArgs(m_tmp_indices, m_bwt_data_sizes, STAGE1_DATA_STRIDE * 8, m_salsa20_keys);
+    m_filter_kernel->setArgs(*job.nonce(), BWT_DATA_MAX_SIZE, m_salsa20_keys, m_filtered_hashes);
+    m_prepare_batch2_kernel->setArgs(m_salsa20_keys, m_filtered_hashes, m_bwt_data_sizes);
+    m_find_shares_kernel->setArgs(m_salsa20_keys, m_filtered_hashes, m_output);
+    m_find_shares_kernel->setTarget(job.target());
+
+    const uint32_t zero = 0;
+    enqueueWriteBuffer(m_filtered_hashes, CL_TRUE, 0, sizeof(uint32_t), &zero);
+}
+
+
+void xmrig::OclAstroBWTRunner::build()
+{
+    OclBaseRunner::build();
+
+    m_sha3_initial_kernel   = new AstroBWT_SHA3InitialKernel(m_program);
+    m_sha3_kernel           = new AstroBWT_SHA3Kernel(m_program);
+    m_salsa20_kernel        = new AstroBWT_Salsa20Kernel(m_program);
+    m_bwt_kernel            = new AstroBWT_MainKernel(m_program);
+    m_filter_kernel         = new AstroBWT_FilterKernel(m_program);
+    m_prepare_batch2_kernel = new AstroBWT_PrepareBatch2Kernel(m_program);
+    m_find_shares_kernel    = new AstroBWT_FindSharesKernel(m_program);
+}
+
+
+void xmrig::OclAstroBWTRunner::init()
+{
+    OclBaseRunner::init();
+
+    const cl_mem_flags f = CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS;
+
+    m_salsa20_keys    = createSubBuffer(f, m_batch_size1 * 32);
+    m_bwt_data        = createSubBuffer(f, m_bwt_allocation_size);
+    m_bwt_data_sizes  = createSubBuffer(CL_MEM_READ_WRITE | CL_MEM_HOST_WRITE_ONLY, m_batch_size1 * 4);
+    m_indices         = createSubBuffer(f, m_bwt_allocation_size * 8);
+    m_tmp_indices     = createSubBuffer(f, m_bwt_allocation_size * 8);
+    m_filtered_hashes = createSubBuffer(CL_MEM_READ_WRITE, (m_batch_size1 + m_intensity) * 36 + 4);
+}
diff --git a/src/backend/opencl/runners/OclAstroBWTRunner.h b/src/backend/opencl/runners/OclAstroBWTRunner.h
new file mode 100644
index 000000000..f16e00017
--- /dev/null
+++ b/src/backend/opencl/runners/OclAstroBWTRunner.h
@@ -0,0 +1,95 @@
+/* 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-2020 SChernykh   <https://github.com/SChernykh>
+ * Copyright 2016-2020 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_OCLASTROBWTRUNNER_H
+#define XMRIG_OCLASTROBWTRUNNER_H
+
+
+#include "backend/opencl/runners/OclBaseRunner.h"
+
+namespace xmrig {
+
+
+class AstroBWT_FilterKernel;
+class AstroBWT_MainKernel;
+class AstroBWT_PrepareBatch2Kernel;
+class AstroBWT_Salsa20Kernel;
+class AstroBWT_SHA3InitialKernel;
+class AstroBWT_SHA3Kernel;
+class AstroBWT_FindSharesKernel;
+
+
+class OclAstroBWTRunner : public OclBaseRunner
+{
+public:
+    static constexpr uint32_t BWT_DATA_MAX_SIZE = 600 * 1024 - 256;
+    static constexpr uint32_t BWT_DATA_STRIDE = (BWT_DATA_MAX_SIZE + 256 + 255) & ~255U;
+
+    XMRIG_DISABLE_COPY_MOVE_DEFAULT(OclAstroBWTRunner)
+
+    OclAstroBWTRunner(size_t index, const OclLaunchData &data);
+    ~OclAstroBWTRunner() override;
+
+    inline uint32_t roundSize() const override { return m_batch_size1; }
+
+    // ~0.5% of all hashes are invalid
+    inline uint32_t processedHashes() const override { return static_cast<uint32_t>(m_processedHashes * 0.995); }
+
+protected:
+    size_t bufferSize() const override;
+    void run(uint32_t nonce, uint32_t *hashOutput) override;
+    void set(const Job &job, uint8_t *blob) override;
+    void build() override;
+    void init() override;
+
+private:
+    AstroBWT_SHA3InitialKernel*   m_sha3_initial_kernel   = nullptr;
+    AstroBWT_SHA3Kernel*          m_sha3_kernel           = nullptr;
+    AstroBWT_Salsa20Kernel*       m_salsa20_kernel        = nullptr;
+    AstroBWT_MainKernel*          m_bwt_kernel            = nullptr;
+    AstroBWT_FilterKernel*        m_filter_kernel         = nullptr;
+    AstroBWT_PrepareBatch2Kernel* m_prepare_batch2_kernel = nullptr;
+    AstroBWT_FindSharesKernel*    m_find_shares_kernel    = nullptr;
+
+
+    cl_mem m_salsa20_keys                                 = nullptr;
+    cl_mem m_bwt_data                                     = nullptr;
+    cl_mem m_bwt_data_sizes                               = nullptr;
+    cl_mem m_indices                                      = nullptr;
+    cl_mem m_tmp_indices                                  = nullptr;
+    cl_mem m_filtered_hashes                              = nullptr;
+
+    uint32_t m_workgroup_size                             = 0;
+    uint64_t m_bwt_allocation_size                        = 0;
+    uint64_t m_batch_size1                                = 0;
+    uint32_t m_processedHashes                                = 0;
+
+    uint32_t* m_bwt_data_sizes_host                       = nullptr;
+};
+
+
+} /* namespace xmrig */
+
+
+#endif // XMRIG_OCLASTROBWTRUNNER_H
diff --git a/src/backend/opencl/runners/OclBaseRunner.h b/src/backend/opencl/runners/OclBaseRunner.h
index 68a20f251..db9f69fd5 100644
--- a/src/backend/opencl/runners/OclBaseRunner.h
+++ b/src/backend/opencl/runners/OclBaseRunner.h
@@ -57,6 +57,8 @@ protected:
     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; }
+    inline uint32_t roundSize() const override          { return m_intensity; }
+    inline uint32_t processedHashes() const override    { return m_intensity; }
 
     size_t bufferSize() const override;
     uint32_t deviceIndex() const override;
diff --git a/src/backend/opencl/wrappers/OclDevice.cpp b/src/backend/opencl/wrappers/OclDevice.cpp
index bd51d0f3f..cd90e95b7 100644
--- a/src/backend/opencl/wrappers/OclDevice.cpp
+++ b/src/backend/opencl/wrappers/OclDevice.cpp
@@ -52,6 +52,10 @@ namespace xmrig {
 extern bool ocl_generic_rx_generator(const OclDevice &device, const Algorithm &algorithm, OclThreads &threads);
 #endif
 
+#ifdef XMRIG_ALGO_ASTROBWT
+extern bool ocl_generic_astrobwt_generator(const OclDevice& device, const Algorithm& algorithm, OclThreads& threads);
+#endif
+
 #ifdef XMRIG_ALGO_CN_GPU
 extern bool ocl_generic_cn_gpu_generator(const OclDevice &device, const Algorithm &algorithm, OclThreads &threads);
 #endif
@@ -64,6 +68,9 @@ static ocl_gen_config_fun generators[] = {
 #   ifdef XMRIG_ALGO_RANDOMX
     ocl_generic_rx_generator,
 #   endif
+#   ifdef XMRIG_ALGO_ASTROBWT
+    ocl_generic_astrobwt_generator,
+#   endif
 #   ifdef XMRIG_ALGO_CN_GPU
     ocl_generic_cn_gpu_generator,
 #   endif