Merge 862280f28c into 77e2f3a028

Fixed Zephyr mining (OpenCL)

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

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

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

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

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

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

src/backend/opencl/opencl.cmake

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

src/backend/opencl/runners/OclRxBaseRunner.cpp

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

src/backend/opencl/runners/OclRxBaseRunner.h

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

src/crypto/cn/CnAlgo.h

@@ -37,14 +37,36 @@ class CnAlgo
 public:
     constexpr CnAlgo() {};
 
-    constexpr inline Algorithm::Id base() const  { static_assert(Algorithm::isCN(ALGO), "invalid CRYPTONIGHT algorithm"); return Algorithm::base(ALGO); }
-    constexpr inline bool isHeavy() const        { return Algorithm::family(ALGO) == Algorithm::CN_HEAVY; }
-    constexpr inline bool isR() const            { return ALGO == Algorithm::CN_R; }
-    constexpr inline size_t memory() const       { static_assert(Algorithm::isCN(ALGO), "invalid CRYPTONIGHT algorithm"); return Algorithm::l3(ALGO); }
-    constexpr inline uint32_t iterations() const { static_assert(Algorithm::isCN(ALGO), "invalid CRYPTONIGHT algorithm"); return CN_ITER; }
+#   define ASSERT_CN static_assert(Algorithm::isCN(ALGO), "invalid CRYPTONIGHT algorithm")
+    constexpr inline Algorithm::Id base() const  { ASSERT_CN; return Algorithm::base(ALGO); }
+    constexpr inline size_t memory() const       { ASSERT_CN; return Algorithm::l3(ALGO); }
+    constexpr inline uint32_t iterations() const { ASSERT_CN; return CN_ITER; }
     constexpr inline uint32_t mask() const       { return static_cast<uint32_t>(((memory() - 1) / 16) * 16); }
     constexpr inline uint32_t half_mem() const   { return mask() < memory() / 2; }
 
+    constexpr inline bool isBase1() const        { ASSERT_CN; return Algorithm::base(ALGO) == Algorithm::CN_1; }
+    constexpr inline bool isBase2() const        { ASSERT_CN; return Algorithm::base(ALGO) == Algorithm::CN_2; }
+    constexpr inline bool is2() const            { return ALGO == Algorithm::CN_2; }
+    constexpr inline bool isR() const            { return ALGO == Algorithm::CN_R; }
+    constexpr inline bool isHalf() const         { return ALGO == Algorithm::CN_HALF; }
+    constexpr inline bool isRTO() const          { return ALGO == Algorithm::CN_RTO; }
+    constexpr inline bool isRWZ() const          { return ALGO == Algorithm::CN_RWZ; }
+    constexpr inline bool isZLS() const          { return ALGO == Algorithm::CN_ZLS; }
+    constexpr inline bool isDouble() const       { return ALGO == Algorithm::CN_DOUBLE; }
+    constexpr inline bool isCCX() const          { return ALGO == Algorithm::CN_CCX; }
+    constexpr inline bool isHeavy() const        { ASSERT_CN; return Algorithm::family(ALGO) == Algorithm::CN_HEAVY; }
+    constexpr inline bool isHeavyTube() const    { return ALGO == Algorithm::CN_HEAVY_TUBE; }
+    constexpr inline bool isHeavyXHV() const     { return ALGO == Algorithm::CN_HEAVY_XHV; }
+    constexpr inline bool isPico0() const        { return ALGO == Algorithm::CN_PICO_0; }
+    constexpr inline bool isPicoTLO() const      { return ALGO == Algorithm::CN_PICO_TLO; }
+    constexpr inline bool isUPX2() const         { return ALGO == Algorithm::CN_UPX2; }
+    constexpr inline bool isGR0() const          { return ALGO == Algorithm::CN_GR_0; }
+    constexpr inline bool isGR1() const          { return ALGO == Algorithm::CN_GR_1; }
+    constexpr inline bool isGR2() const          { return ALGO == Algorithm::CN_GR_2; }
+    constexpr inline bool isGR3() const          { return ALGO == Algorithm::CN_GR_3; }
+    constexpr inline bool isGR4() const          { return ALGO == Algorithm::CN_GR_4; }
+    constexpr inline bool isGR5() const          { return ALGO == Algorithm::CN_GR_5; }
+
     inline static uint32_t iterations(Algorithm::Id algo)
     {
         switch (algo) {

src/crypto/cn/CryptoNight_x86.h

@@ -603,7 +603,7 @@ static inline void cryptonight_monero_tweak(uint64_t *mem_out, const uint8_t *l,
     constexpr CnAlgo<ALGO> props;
 
     if (props.base() == Algorithm::CN_2) {
-        VARIANT2_SHUFFLE(l, idx, ax0, bx0, bx1, cx, (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? 1 : 0));
+        VARIANT2_SHUFFLE(l, idx, ax0, bx0, bx1, cx, ((props.isRWZ() || props.isUPX2()) ? 1 : 0));
         _mm_store_si128(reinterpret_cast<__m128i *>(mem_out), _mm_xor_si128(bx0, cx));
     } else {
         __m128i tmp = _mm_xor_si128(bx0, cx);
@@ -665,15 +665,8 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
 
     constexpr CnAlgo<ALGO> props;
     constexpr size_t MASK        = props.mask();
-    constexpr Algorithm::Id BASE = props.base();
 
-#   ifdef XMRIG_ALGO_CN_HEAVY
-    constexpr bool IS_CN_HEAVY_TUBE = ALGO == Algorithm::CN_HEAVY_TUBE;
-#   else
-    constexpr bool IS_CN_HEAVY_TUBE = false;
-#   endif
-
-    if (BASE == Algorithm::CN_1 && size < 43) {
+    if (props.isBase1() && size < 43) {
         memset(output, 0, 32);
         return;
     }
@@ -694,10 +687,7 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
             V4_Instruction code[256];
             const int code_size = v4_random_math_init<ALGO>(code, height);
 
-            if (ALGO == Algorithm::CN_R) {
-                v4_soft_aes_compile_code(code, code_size, reinterpret_cast<void*>(ctx[0]->generated_code), Assembly::NONE);
-            }
-
+            v4_soft_aes_compile_code(code, code_size, reinterpret_cast<void*>(ctx[0]->generated_code), Assembly::NONE);
             ctx[0]->generated_code_data = { ALGO, height };
         }
 
@@ -718,26 +708,26 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
     __m128i bx1   = _mm_set_epi64x(static_cast<int64_t>(h0[9] ^ h0[11]), static_cast<int64_t>(h0[8] ^ h0[10]));
 
     __m128 conc_var;
-    if (ALGO == Algorithm::CN_CCX) {
+    if (props.isCCX()) {
         conc_var = _mm_setzero_ps();
         RESTORE_ROUNDING_MODE();
     }
 
     for (size_t i = 0; i < props.iterations(); i++) {
         __m128i cx;
-        if (IS_CN_HEAVY_TUBE || !SOFT_AES) {
+        if (props.isHeavyTube() || !SOFT_AES) {
             cx = _mm_load_si128(reinterpret_cast<const __m128i *>(&l0[interleaved_index<interleave>(idx0 & MASK)]));
-            if (ALGO == Algorithm::CN_CCX) {
+            if (props.isCCX()) {
                 cryptonight_conceal_tweak(cx, conc_var);
             }
         }
 
         const __m128i ax0 = _mm_set_epi64x(static_cast<int64_t>(ah0), static_cast<int64_t>(al0));
-        if (IS_CN_HEAVY_TUBE) {
+        if (props.isHeavyTube()) {
             cx = aes_round_tweak_div(cx, ax0);
         }
         else if (SOFT_AES) {
-            if (ALGO == Algorithm::CN_CCX) {
+            if (props.isCCX()) {
                 cx = _mm_load_si128(reinterpret_cast<const __m128i*>(&l0[interleaved_index<interleave>(idx0 & MASK)]));
                 cryptonight_conceal_tweak(cx, conc_var);
                 cx = soft_aesenc(&cx, ax0, reinterpret_cast<const uint32_t*>(saes_table));
@@ -750,7 +740,7 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
             cx = _mm_aesenc_si128(cx, ax0);
         }
 
-        if (BASE == Algorithm::CN_1 || BASE == Algorithm::CN_2) {
+        if (props.isBase1() || props.isBase2()) {
             cryptonight_monero_tweak<ALGO>(reinterpret_cast<uint64_t*>(&l0[interleaved_index<interleave>(idx0 & MASK)]), l0, idx0 & MASK, ax0, bx0, bx1, cx);
         } else {
             _mm_store_si128(reinterpret_cast<__m128i *>(&l0[interleaved_index<interleave>(idx0 & MASK)]), _mm_xor_si128(bx0, cx));
@@ -762,13 +752,11 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
         cl = (reinterpret_cast<uint64_t*>(&l0[interleaved_index<interleave>(idx0 & MASK)]))[0];
         ch = (reinterpret_cast<uint64_t*>(&l0[interleaved_index<interleave>(idx0 & MASK)]))[1];
 
-        if (BASE == Algorithm::CN_2) {
+        if (props.isBase2()) {
             if (props.isR()) {
                 VARIANT4_RANDOM_MATH(0, al0, ah0, cl, bx0, bx1);
-                if (ALGO == Algorithm::CN_R) {
-                    al0 ^= r0[2] | (static_cast<uint64_t>(r0[3]) << 32);
-                    ah0 ^= r0[0] | (static_cast<uint64_t>(r0[1]) << 32);
-                }
+                al0 ^= r0[2] | (static_cast<uint64_t>(r0[3]) << 32);
+                ah0 ^= r0[0] | (static_cast<uint64_t>(r0[1]) << 32);
             } else {
                 VARIANT2_INTEGER_MATH(0, cl, cx);
             }
@@ -776,11 +764,11 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
 
         lo = __umul128(idx0, cl, &hi);
 
-        if (BASE == Algorithm::CN_2) {
-            if (ALGO == Algorithm::CN_R) {
+        if (props.isBase2()) {
+            if (props.isR()) {
                 VARIANT2_SHUFFLE(l0, idx0 & MASK, ax0, bx0, bx1, cx, 0);
             } else {
-                VARIANT2_SHUFFLE2(l0, idx0 & MASK, ax0, bx0, bx1, hi, lo, (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? 1 : 0));
+                VARIANT2_SHUFFLE2(l0, idx0 & MASK, ax0, bx0, bx1, hi, lo, ((props.isRWZ() || props.isUPX2()) ? 1 : 0));
             }
         }
 
@@ -789,9 +777,9 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
 
         reinterpret_cast<uint64_t*>(&l0[interleaved_index<interleave>(idx0 & MASK)])[0] = al0;
 
-        if (IS_CN_HEAVY_TUBE || ALGO == Algorithm::CN_RTO) {
+        if (props.isHeavyTube() || props.isRTO()) {
             reinterpret_cast<uint64_t*>(&l0[interleaved_index<interleave>(idx0 & MASK)])[1] = ah0 ^ tweak1_2_0 ^ al0;
-        } else if (BASE == Algorithm::CN_1) {
+        } else if (props.isBase1()) {
             reinterpret_cast<uint64_t*>(&l0[interleaved_index<interleave>(idx0 & MASK)])[1] = ah0 ^ tweak1_2_0;
         } else {
             reinterpret_cast<uint64_t*>(&l0[interleaved_index<interleave>(idx0 & MASK)])[1] = ah0;
@@ -819,7 +807,7 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
 
             ((int64_t*)&l0[interleaved_index<interleave>(idx0 & MASK)])[0] = n ^ q;
 
-            if (ALGO == Algorithm::CN_HEAVY_XHV) {
+            if (props.isHeavyXHV()) {
                 d = ~d;
             }
 
@@ -827,7 +815,7 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
         }
 #       endif
 
-        if (BASE == Algorithm::CN_2) {
+        if (props.isBase2()) {
             bx1 = bx0;
         }
 
@@ -960,7 +948,7 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
     }
     cn_explode_scratchpad<ALGO, false, 0>(ctx[0]);
 
-    if (ALGO == Algorithm::CN_2) {
+    if (props.is2()) {
         if (ASM == Assembly::INTEL) {
             cnv2_mainloop_ivybridge_asm(ctx);
         }
@@ -971,7 +959,7 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
             cnv2_mainloop_bulldozer_asm(ctx);
         }
     }
-    else if (ALGO == Algorithm::CN_HALF) {
+    else if (props.isHalf()) {
         if (ASM == Assembly::INTEL) {
             cn_half_mainloop_ivybridge_asm(ctx);
         }
@@ -983,7 +971,7 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
         }
     }
 #   ifdef XMRIG_ALGO_CN_PICO
-    else if (ALGO == Algorithm::CN_PICO_0) {
+    else if (props.isPico0()) {
         if (ASM == Assembly::INTEL) {
             cn_trtl_mainloop_ivybridge_asm(ctx);
         }
@@ -994,7 +982,7 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
             cn_trtl_mainloop_bulldozer_asm(ctx);
         }
     }
-    else if (ALGO == Algorithm::CN_PICO_TLO) {
+    else if (props.isPicoTLO()) {
         if (ASM == Assembly::INTEL) {
             cn_tlo_mainloop_ivybridge_asm(ctx);
         }
@@ -1006,10 +994,10 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
         }
     }
 #   endif
-    else if (ALGO == Algorithm::CN_RWZ) {
+    else if (props.isRWZ()) {
         cnv2_rwz_mainloop_asm(ctx);
     }
-    else if (ALGO == Algorithm::CN_ZLS) {
+    else if (props.isZLS()) {
         if (ASM == Assembly::INTEL) {
             cn_zls_mainloop_ivybridge_asm(ctx);
         }
@@ -1020,7 +1008,7 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
             cn_zls_mainloop_bulldozer_asm(ctx);
         }
     }
-    else if (ALGO == Algorithm::CN_DOUBLE) {
+    else if (props.isDouble()) {
         if (ASM == Assembly::INTEL) {
             cn_double_mainloop_ivybridge_asm(ctx);
         }
@@ -1032,7 +1020,7 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
         }
     }
 #   ifdef XMRIG_ALGO_CN_FEMTO
-    else if (ALGO == Algorithm::CN_UPX2) {
+    else if (props.isUPX2()) {
         cn_upx2_mainloop_asm(ctx);
     }
 #   endif
@@ -1078,22 +1066,22 @@ inline void cryptonight_double_hash_asm(const uint8_t *__restrict__ input, size_
         cn_explode_scratchpad<ALGO, false, 0>(ctx[1]);
     }
 
-    if (ALGO == Algorithm::CN_2) {
+    if (props.is2()) {
         cnv2_double_mainloop_sandybridge_asm(ctx);
     }
-    else if (ALGO == Algorithm::CN_HALF) {
+    else if (props.isHalf()){
         cn_half_double_mainloop_sandybridge_asm(ctx);
     }
 #   ifdef XMRIG_ALGO_CN_PICO
-    else if (ALGO == Algorithm::CN_PICO_0) {
+    else if (props.isPico0()) {
         cn_trtl_double_mainloop_sandybridge_asm(ctx);
     }
-    else if (ALGO == Algorithm::CN_PICO_TLO) {
+    else if (props.isPicoTLO()) {
         cn_tlo_double_mainloop_sandybridge_asm(ctx);
     }
 #   endif
 #   ifdef XMRIG_ALGO_CN_FEMTO
-    else if (ALGO == Algorithm::CN_UPX2) {
+    else if (props.isUPX2()) {
         if (Cpu::info()->arch() == ICpuInfo::ARCH_ZEN3) {
             cnv2_upx_double_mainloop_zen3_asm(ctx);
         }
@@ -1102,13 +1090,13 @@ inline void cryptonight_double_hash_asm(const uint8_t *__restrict__ input, size_
         }
     }
 #   endif
-    else if (ALGO == Algorithm::CN_RWZ) {
+    else if (props.isRWZ()) {
         cnv2_rwz_double_mainloop_asm(ctx);
     }
-    else if (ALGO == Algorithm::CN_ZLS) {
+    else if (props.isZLS()) {
         cn_zls_double_mainloop_sandybridge_asm(ctx);
     }
-    else if (ALGO == Algorithm::CN_DOUBLE) {
+    else if (props.isDouble()) {
         cn_double_double_mainloop_sandybridge_asm(ctx);
     }
     else if (props.isR()) {
@@ -1146,9 +1134,8 @@ template<Algorithm::Id ALGO>
 static NOINLINE void cryptonight_single_hash_gr_sse41(const uint8_t* __restrict__ input, size_t size, uint8_t* __restrict__ output, cryptonight_ctx** __restrict__ ctx, uint64_t height)
 {
     constexpr CnAlgo<ALGO> props;
-    constexpr Algorithm::Id BASE = props.base();
 
-    if (BASE == Algorithm::CN_1 && size < 43) {
+    if (props.isBase1() && size < 43) {
         memset(output, 0, 32);
         return;
     }
@@ -1163,12 +1150,12 @@ static NOINLINE void cryptonight_single_hash_gr_sse41(const uint8_t* __restrict_
     VARIANT1_INIT(0);
     ctx[0]->tweak1_2 = tweak1_2_0;
     ctx[0]->tweak1_table = tweak1_table;
-    if (ALGO == Algorithm::CN_GR_0) cn_gr0_single_mainloop_asm(ctx);
-    if (ALGO == Algorithm::CN_GR_1) cn_gr1_single_mainloop_asm(ctx);
-    if (ALGO == Algorithm::CN_GR_2) cn_gr2_single_mainloop_asm(ctx);
-    if (ALGO == Algorithm::CN_GR_3) cn_gr3_single_mainloop_asm(ctx);
-    if (ALGO == Algorithm::CN_GR_4) cn_gr4_single_mainloop_asm(ctx);
-    if (ALGO == Algorithm::CN_GR_5) cn_gr5_single_mainloop_asm(ctx);
+    if (props.isGR0()) cn_gr0_single_mainloop_asm(ctx);
+    if (props.isGR1()) cn_gr1_single_mainloop_asm(ctx);
+    if (props.isGR2()) cn_gr2_single_mainloop_asm(ctx);
+    if (props.isGR3()) cn_gr3_single_mainloop_asm(ctx);
+    if (props.isGR4()) cn_gr4_single_mainloop_asm(ctx);
+    if (props.isGR5()) cn_gr5_single_mainloop_asm(ctx);
 
     cn_implode_scratchpad<ALGO, false, 0>(ctx[0]);
     keccakf(reinterpret_cast<uint64_t*>(ctx[0]->state), 24);
@@ -1180,9 +1167,8 @@ template<Algorithm::Id ALGO>
 static NOINLINE void cryptonight_double_hash_gr_sse41(const uint8_t *__restrict__ input, size_t size, uint8_t *__restrict__ output, cryptonight_ctx **__restrict__ ctx, uint64_t height)
 {
     constexpr CnAlgo<ALGO> props;
-    constexpr Algorithm::Id BASE = props.base();
 
-    if (BASE == Algorithm::CN_1 && size < 43) {
+    if (props.isBase1() && size < 43) {
         memset(output, 0, 64);
         return;
     }
@@ -1196,7 +1182,7 @@ static NOINLINE void cryptonight_double_hash_gr_sse41(const uint8_t *__restrict_
     }
 
 #   ifdef XMRIG_VAES
-    if (!props.isHeavy() && cn_vaes_enabled) {
+    if (cn_vaes_enabled) {
         cn_explode_scratchpad_vaes_double(ctx[0], ctx[1], props.memory(), props.half_mem());
     }
     else
@@ -1214,15 +1200,15 @@ static NOINLINE void cryptonight_double_hash_gr_sse41(const uint8_t *__restrict_
 
     ctx[0]->tweak1_table = tweak1_table;
 
-    if (ALGO == Algorithm::CN_GR_0) cn_gr0_double_mainloop_asm(ctx);
-    if (ALGO == Algorithm::CN_GR_1) cn_gr1_double_mainloop_asm(ctx);
-    if (ALGO == Algorithm::CN_GR_2) cn_gr2_double_mainloop_asm(ctx);
-    if (ALGO == Algorithm::CN_GR_3) cn_gr3_double_mainloop_asm(ctx);
-    if (ALGO == Algorithm::CN_GR_4) cn_gr4_double_mainloop_asm(ctx);
-    if (ALGO == Algorithm::CN_GR_5) cn_gr5_double_mainloop_asm(ctx);
+    if (props.isGR0()) cn_gr0_double_mainloop_asm(ctx);
+    if (props.isGR1()) cn_gr1_double_mainloop_asm(ctx);
+    if (props.isGR2()) cn_gr2_double_mainloop_asm(ctx);
+    if (props.isGR3()) cn_gr3_double_mainloop_asm(ctx);
+    if (props.isGR4()) cn_gr4_double_mainloop_asm(ctx);
+    if (props.isGR5()) cn_gr5_double_mainloop_asm(ctx);
 
 #   ifdef XMRIG_VAES
-    if (!props.isHeavy() && cn_vaes_enabled) {
+    if (cn_vaes_enabled) {
         cn_implode_scratchpad_vaes_double(ctx[0], ctx[1], props.memory(), props.half_mem());
     }
     else
@@ -1267,15 +1253,8 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
 
     constexpr CnAlgo<ALGO> props;
     constexpr size_t MASK        = props.mask();
-    constexpr Algorithm::Id BASE = props.base();
 
-#   ifdef XMRIG_ALGO_CN_HEAVY
-    constexpr bool IS_CN_HEAVY_TUBE = ALGO == Algorithm::CN_HEAVY_TUBE;
-#   else
-    constexpr bool IS_CN_HEAVY_TUBE = false;
-#   endif
-
-    if (BASE == Algorithm::CN_1 && size < 43) {
+    if (props.isBase1() && size < 43) {
         memset(output, 0, 64);
         return;
     }
@@ -1323,7 +1302,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
     __m128i bx11 = _mm_set_epi64x(h1[9] ^ h1[11], h1[8] ^ h1[10]);
 
     __m128 conc_var0, conc_var1;
-    if (ALGO == Algorithm::CN_CCX) {
+    if (props.isCCX()) {
         conc_var0 = _mm_setzero_ps();
         conc_var1 = _mm_setzero_ps();
         RESTORE_ROUNDING_MODE();
@@ -1334,10 +1313,10 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
 
     for (size_t i = 0; i < props.iterations(); i++) {
         __m128i cx0, cx1;
-        if (IS_CN_HEAVY_TUBE || !SOFT_AES) {
+        if (props.isHeavyTube() || !SOFT_AES) {
             cx0 = _mm_load_si128(reinterpret_cast<const __m128i *>(&l0[idx0 & MASK]));
             cx1 = _mm_load_si128(reinterpret_cast<const __m128i *>(&l1[idx1 & MASK]));
-            if (ALGO == Algorithm::CN_CCX) {
+            if (props.isCCX()) {
                 cryptonight_conceal_tweak(cx0, conc_var0);
                 cryptonight_conceal_tweak(cx1, conc_var1);
             }
@@ -1345,12 +1324,12 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
 
         const __m128i ax0 = _mm_set_epi64x(ah0, al0);
         const __m128i ax1 = _mm_set_epi64x(ah1, al1);
-        if (IS_CN_HEAVY_TUBE) {
+        if (props.isHeavyTube()) {
             cx0 = aes_round_tweak_div(cx0, ax0);
             cx1 = aes_round_tweak_div(cx1, ax1);
         }
         else if (SOFT_AES) {
-            if (ALGO == Algorithm::CN_CCX) {
+            if (props.isCCX()) {
                 cx0 = _mm_load_si128(reinterpret_cast<const __m128i*>(&l0[idx0 & MASK]));
                 cx1 = _mm_load_si128(reinterpret_cast<const __m128i*>(&l1[idx1 & MASK]));
                 cryptonight_conceal_tweak(cx0, conc_var0);
@@ -1368,7 +1347,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
             cx1 = _mm_aesenc_si128(cx1, ax1);
         }
 
-        if (BASE == Algorithm::CN_1 || BASE == Algorithm::CN_2) {
+        if (props.isBase1() || props.isBase2()) {
             cryptonight_monero_tweak<ALGO>((uint64_t*)&l0[idx0 & MASK], l0, idx0 & MASK, ax0, bx00, bx01, cx0);
             cryptonight_monero_tweak<ALGO>((uint64_t*)&l1[idx1 & MASK], l1, idx1 & MASK, ax1, bx10, bx11, cx1);
         } else {
@@ -1383,13 +1362,11 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
         cl = ((uint64_t*) &l0[idx0 & MASK])[0];
         ch = ((uint64_t*) &l0[idx0 & MASK])[1];
 
-        if (BASE == Algorithm::CN_2) {
+        if (props.isBase2()) {
             if (props.isR()) {
                 VARIANT4_RANDOM_MATH(0, al0, ah0, cl, bx00, bx01);
-                if (ALGO == Algorithm::CN_R) {
-                    al0 ^= r0[2] | ((uint64_t)(r0[3]) << 32);
-                    ah0 ^= r0[0] | ((uint64_t)(r0[1]) << 32);
-                }
+                al0 ^= r0[2] | ((uint64_t)(r0[3]) << 32);
+                ah0 ^= r0[0] | ((uint64_t)(r0[1]) << 32);
             } else {
                 VARIANT2_INTEGER_MATH(0, cl, cx0);
             }
@@ -1397,11 +1374,11 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
 
         lo = __umul128(idx0, cl, &hi);
 
-        if (BASE == Algorithm::CN_2) {
-            if (ALGO == Algorithm::CN_R) {
+        if (props.isBase2()) {
+            if (props.isR()) {
                 VARIANT2_SHUFFLE(l0, idx0 & MASK, ax0, bx00, bx01, cx0, 0);
             } else {
-                VARIANT2_SHUFFLE2(l0, idx0 & MASK, ax0, bx00, bx01, hi, lo, (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? 1 : 0));
+                VARIANT2_SHUFFLE2(l0, idx0 & MASK, ax0, bx00, bx01, hi, lo, ((props.isRWZ() || props.isUPX2()) ? 1 : 0));
             }
         }
 
@@ -1410,9 +1387,9 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
 
         ((uint64_t*)&l0[idx0 & MASK])[0] = al0;
 
-        if (IS_CN_HEAVY_TUBE || ALGO == Algorithm::CN_RTO) {
+        if (props.isHeavyTube() || props.isRTO()) {
             ((uint64_t*) &l0[idx0 & MASK])[1] = ah0 ^ tweak1_2_0 ^ al0;
-        } else if (BASE == Algorithm::CN_1) {
+        } else if (props.isBase1()) {
             ((uint64_t*) &l0[idx0 & MASK])[1] = ah0 ^ tweak1_2_0;
         } else {
             ((uint64_t*) &l0[idx0 & MASK])[1] = ah0;
@@ -1430,7 +1407,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
 
             ((int64_t*)&l0[idx0 & MASK])[0] = n ^ q;
 
-            if (ALGO == Algorithm::CN_HEAVY_XHV) {
+            if (props.isHeavyXHV()) {
                 d = ~d;
             }
 
@@ -1441,13 +1418,11 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
         cl = ((uint64_t*) &l1[idx1 & MASK])[0];
         ch = ((uint64_t*) &l1[idx1 & MASK])[1];
 
-        if (BASE == Algorithm::CN_2) {
+        if (props.isBase2()) {
             if (props.isR()) {
                 VARIANT4_RANDOM_MATH(1, al1, ah1, cl, bx10, bx11);
-                if (ALGO == Algorithm::CN_R) {
-                    al1 ^= r1[2] | ((uint64_t)(r1[3]) << 32);
-                    ah1 ^= r1[0] | ((uint64_t)(r1[1]) << 32);
-                }
+                al1 ^= r1[2] | ((uint64_t)(r1[3]) << 32);
+                ah1 ^= r1[0] | ((uint64_t)(r1[1]) << 32);
             } else {
                 VARIANT2_INTEGER_MATH(1, cl, cx1);
             }
@@ -1455,11 +1430,11 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
 
         lo = __umul128(idx1, cl, &hi);
 
-        if (BASE == Algorithm::CN_2) {
-            if (ALGO == Algorithm::CN_R) {
+        if (props.isBase2()) {
+            if (props.isR()) {
                 VARIANT2_SHUFFLE(l1, idx1 & MASK, ax1, bx10, bx11, cx1, 0);
             } else {
-                VARIANT2_SHUFFLE2(l1, idx1 & MASK, ax1, bx10, bx11, hi, lo, (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? 1 : 0));
+                VARIANT2_SHUFFLE2(l1, idx1 & MASK, ax1, bx10, bx11, hi, lo, ((props.isRWZ() || props.isUPX2()) ? 1 : 0));
             }
         }
 
@@ -1468,9 +1443,9 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
 
         ((uint64_t*)&l1[idx1 & MASK])[0] = al1;
 
-        if (IS_CN_HEAVY_TUBE || ALGO == Algorithm::CN_RTO) {
+        if (props.isHeavyTube() || props.isRTO()) {
             ((uint64_t*)&l1[idx1 & MASK])[1] = ah1 ^ tweak1_2_1 ^ al1;
-        } else if (BASE == Algorithm::CN_1) {
+        } else if (props.isBase1()) {
             ((uint64_t*)&l1[idx1 & MASK])[1] = ah1 ^ tweak1_2_1;
         } else {
             ((uint64_t*)&l1[idx1 & MASK])[1] = ah1;
@@ -1488,7 +1463,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
 
             ((int64_t*)&l1[idx1 & MASK])[0] = n ^ q;
 
-            if (ALGO == Algorithm::CN_HEAVY_XHV) {
+            if (props.isHeavyXHV()) {
                 d = ~d;
             }
 
@@ -1496,7 +1471,7 @@ inline void cryptonight_double_hash(const uint8_t *__restrict__ input, size_t si
         }
 #       endif
 
-        if (BASE == Algorithm::CN_2) {
+        if (props.isBase2()) {
             bx01 = bx00;
             bx11 = bx10;
         }
@@ -1529,9 +1504,8 @@ template<Algorithm::Id ALGO>
 static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__ input, size_t size, uint8_t* __restrict__ output, cryptonight_ctx** __restrict__ ctx, uint64_t height)
 {
     constexpr CnAlgo<ALGO> props;
-    constexpr Algorithm::Id BASE = props.base();
 
-    if (BASE == Algorithm::CN_1 && size < 43) {
+    if (props.isBase1() && size < 43) {
         memset(output, 0, 32 * 4);
         return;
     }
@@ -1549,7 +1523,7 @@ static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__
     }
 
 #   ifdef XMRIG_VAES
-    if (!props.isHeavy() && cn_vaes_enabled) {
+    if (cn_vaes_enabled) {
         cn_explode_scratchpad_vaes_double(ctx[0], ctx[1], props.memory(), props.half_mem());
         cn_explode_scratchpad_vaes_double(ctx[2], ctx[3], props.memory(), props.half_mem());
     }
@@ -1569,15 +1543,15 @@ static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__
 
     ctx[0]->tweak1_table = tweak1_table;
 
-    if (ALGO == Algorithm::CN_GR_0) cn_gr0_quad_mainloop_asm(ctx);
-    if (ALGO == Algorithm::CN_GR_1) cn_gr1_quad_mainloop_asm(ctx);
-    if (ALGO == Algorithm::CN_GR_2) cn_gr2_quad_mainloop_asm(ctx);
-    if (ALGO == Algorithm::CN_GR_3) cn_gr3_quad_mainloop_asm(ctx);
-    if (ALGO == Algorithm::CN_GR_4) cn_gr4_quad_mainloop_asm(ctx);
-    if (ALGO == Algorithm::CN_GR_5) cn_gr5_quad_mainloop_asm(ctx);
+    if (props.isGR0()) cn_gr0_quad_mainloop_asm(ctx);
+    if (props.isGR1()) cn_gr1_quad_mainloop_asm(ctx);
+    if (props.isGR2()) cn_gr2_quad_mainloop_asm(ctx);
+    if (props.isGR3()) cn_gr3_quad_mainloop_asm(ctx);
+    if (props.isGR4()) cn_gr4_quad_mainloop_asm(ctx);
+    if (props.isGR5()) cn_gr5_quad_mainloop_asm(ctx);
 
 #   ifdef XMRIG_VAES
-    if (!props.isHeavy() && cn_vaes_enabled) {
+    if (cn_vaes_enabled) {
         cn_implode_scratchpad_vaes_double(ctx[0], ctx[1], props.memory(), props.half_mem());
         cn_implode_scratchpad_vaes_double(ctx[2], ctx[3], props.memory(), props.half_mem());
     }
@@ -1606,14 +1580,14 @@ static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__
 #define CN_STEP1(a, b0, b1, c, l, ptr, idx, conc_var) \
     ptr = reinterpret_cast<__m128i*>(&l[idx & MASK]); \
     c = _mm_load_si128(ptr);                          \
-    if (ALGO == Algorithm::CN_CCX) {                  \
+    if (props.isCCX()) {                              \
         cryptonight_conceal_tweak(c, conc_var);       \
     }
 
 
 
 #define CN_STEP2(a, b0, b1, c, l, ptr, idx)                                             \
-    if (IS_CN_HEAVY_TUBE) {                                                             \
+    if (props.isHeavyTube()) {                                                          \
         c = aes_round_tweak_div(c, a);                                                  \
     }                                                                                   \
     else if (SOFT_AES) {                                                                \
@@ -1622,7 +1596,7 @@ static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__
         c = _mm_aesenc_si128(c, a);                                                     \
     }                                                                                   \
                                                                                         \
-    if (BASE == Algorithm::CN_1 || BASE == Algorithm::CN_2) {                           \
+    if (props.isBase1() || props.isBase2()) {                                           \
         cryptonight_monero_tweak<ALGO>((uint64_t*)ptr, l, idx & MASK, a, b0, b1, c);    \
     } else {                                                                            \
         _mm_store_si128(ptr, _mm_xor_si128(b0, c));                                     \
@@ -1638,36 +1612,34 @@ static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__
 
 #define CN_STEP4(part, a, b0, b1, c, l, mc, ptr, idx)                                                       \
     uint64_t al##part, ah##part;                                                                            \
-    if (BASE == Algorithm::CN_2) {                                                                          \
+    if (props.isBase2()) {                                                                                  \
         if (props.isR()) {                                                                                  \
             al##part = _mm_cvtsi128_si64(a);                                                                \
             ah##part = _mm_cvtsi128_si64(_mm_srli_si128(a, 8));                                             \
             VARIANT4_RANDOM_MATH(part, al##part, ah##part, cl##part, b0, b1);                               \
-            if (ALGO == Algorithm::CN_R) {                                                                  \
-                al##part ^= r##part[2] | ((uint64_t)(r##part[3]) << 32);                                    \
-                ah##part ^= r##part[0] | ((uint64_t)(r##part[1]) << 32);                                    \
-            }                                                                                               \
+            al##part ^= r##part[2] | ((uint64_t)(r##part[3]) << 32);                                        \
+            ah##part ^= r##part[0] | ((uint64_t)(r##part[1]) << 32);                                        \
         } else {                                                                                            \
             VARIANT2_INTEGER_MATH(part, cl##part, c);                                                       \
         }                                                                                                   \
     }                                                                                                       \
     lo = __umul128(idx, cl##part, &hi);                                                                     \
-    if (BASE == Algorithm::CN_2) {                                                                          \
-        if (ALGO == Algorithm::CN_R) {                                                                      \
+    if (props.isBase2()) {                                                                                  \
+        if (props.isR()) {                                                                                  \
             VARIANT2_SHUFFLE(l, idx & MASK, a, b0, b1, c, 0);                                               \
         } else {                                                                                            \
-            VARIANT2_SHUFFLE2(l, idx & MASK, a, b0, b1, hi, lo, (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? 1 : 0)); \
+            VARIANT2_SHUFFLE2(l, idx & MASK, a, b0, b1, hi, lo, ((props.isRWZ() || props.isUPX2()) ? 1 : 0)); \
         }                                                                                                   \
     }                                                                                                       \
-    if (ALGO == Algorithm::CN_R) {                                                                          \
+    if (props.isR()) {                                                                                      \
         a = _mm_set_epi64x(ah##part, al##part);                                                             \
     }                                                                                                       \
     a = _mm_add_epi64(a, _mm_set_epi64x(lo, hi));                                                           \
                                                                                                             \
-    if (BASE == Algorithm::CN_1) {                                                                          \
+    if (props.isBase1()) {                                                                                  \
         _mm_store_si128(ptr, _mm_xor_si128(a, mc));                                                         \
                                                                                                             \
-        if (IS_CN_HEAVY_TUBE || ALGO == Algorithm::CN_RTO) {                                                \
+        if (props.isHeavyTube() || props.isRTO()) {                                                         \
             ((uint64_t*)ptr)[1] ^= ((uint64_t*)ptr)[0];                                                     \
         }                                                                                                   \
     } else {                                                                                                \
@@ -1681,13 +1653,13 @@ static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__
         int32_t d = ((int32_t*)&l[idx & MASK])[2];                                                          \
         int64_t q = n / (d | 0x5);                                                                          \
         ((int64_t*)&l[idx & MASK])[0] = n ^ q;                                                              \
-        if (IS_CN_HEAVY_XHV) {                                                                              \
+        if (props.isHeavyXHV()) {                                                                           \
             d = ~d;                                                                                         \
         }                                                                                                   \
                                                                                                             \
         idx = d ^ q;                                                                                        \
     }                                                                                                       \
-    if (BASE == Algorithm::CN_2) {                                                                          \
+    if (props.isBase2()) {                                                                                  \
         b1 = b0;                                                                                            \
     }                                                                                                       \
     b0 = c;
@@ -1697,11 +1669,11 @@ static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__
     __m128i mc##n;                                                                               \
     __m128i division_result_xmm_##n;                                                             \
     __m128i sqrt_result_xmm_##n;                                                                 \
-    if (BASE == Algorithm::CN_1) {                                                               \
+    if (props.isBase1()) {                                                                       \
         mc##n = _mm_set_epi64x(*reinterpret_cast<const uint64_t*>(input + n * size + 35) ^       \
                                *(reinterpret_cast<const uint64_t*>((ctx)->state) + 24), 0);      \
     }                                                                                            \
-    if (BASE == Algorithm::CN_2) {                                                               \
+    if (props.isBase2()) {                                                                       \
         division_result_xmm_##n = _mm_cvtsi64_si128(h##n[12]);                                   \
         sqrt_result_xmm_##n = _mm_cvtsi64_si128(h##n[13]);                                       \
     }                                                                                            \
@@ -1710,7 +1682,7 @@ static NOINLINE void cryptonight_quad_hash_gr_sse41(const uint8_t* __restrict__
     __m128i bx##n##1 = _mm_set_epi64x(h##n[9] ^ h##n[11], h##n[8] ^ h##n[10]);                   \
     __m128i cx##n = _mm_setzero_si128();                                                         \
     __m128 conc_var##n;                                                                          \
-    if (ALGO == Algorithm::CN_CCX) {                                                             \
+    if (props.isCCX()) {                                                                         \
         conc_var##n = _mm_setzero_ps();                                                          \
     }                                                                                            \
     VARIANT4_RANDOM_MATH_INIT(n);
@@ -1721,17 +1693,8 @@ inline void cryptonight_triple_hash(const uint8_t *__restrict__ input, size_t si
 {
     constexpr CnAlgo<ALGO> props;
     constexpr size_t MASK        = props.mask();
-    constexpr Algorithm::Id BASE = props.base();
 
-#   ifdef XMRIG_ALGO_CN_HEAVY
-    constexpr bool IS_CN_HEAVY_TUBE = ALGO == Algorithm::CN_HEAVY_TUBE;
-    constexpr bool IS_CN_HEAVY_XHV  = ALGO == Algorithm::CN_HEAVY_XHV;
-#   else
-    constexpr bool IS_CN_HEAVY_TUBE = false;
-    constexpr bool IS_CN_HEAVY_XHV  = false;
-#   endif
-
-    if (BASE == Algorithm::CN_1 && size < 43) {
+    if (props.isBase1() && size < 43) {
         memset(output, 0, 32 * 3);
         return;
     }
@@ -1755,7 +1718,7 @@ inline void cryptonight_triple_hash(const uint8_t *__restrict__ input, size_t si
     CONST_INIT(ctx[1], 1);
     CONST_INIT(ctx[2], 2);
     VARIANT2_SET_ROUNDING_MODE();
-    if (ALGO == Algorithm::CN_CCX) {
+    if (props.isCCX()) {
         RESTORE_ROUNDING_MODE();
     }
 
@@ -1819,17 +1782,8 @@ inline void cryptonight_quad_hash(const uint8_t *__restrict__ input, size_t size
 
     constexpr CnAlgo<ALGO> props;
     constexpr size_t MASK        = props.mask();
-    constexpr Algorithm::Id BASE = props.base();
 
-#   ifdef XMRIG_ALGO_CN_HEAVY
-    constexpr bool IS_CN_HEAVY_TUBE = ALGO == Algorithm::CN_HEAVY_TUBE;
-    constexpr bool IS_CN_HEAVY_XHV  = ALGO == Algorithm::CN_HEAVY_XHV;
-#   else
-    constexpr bool IS_CN_HEAVY_TUBE = false;
-    constexpr bool IS_CN_HEAVY_XHV  = false;
-#   endif
-
-    if (BASE == Algorithm::CN_1 && size < 43) {
+    if (props.isBase1() && size < 43) {
         memset(output, 0, 32 * 4);
         return;
     }
@@ -1869,7 +1823,7 @@ inline void cryptonight_quad_hash(const uint8_t *__restrict__ input, size_t size
     CONST_INIT(ctx[2], 2);
     CONST_INIT(ctx[3], 3);
     VARIANT2_SET_ROUNDING_MODE();
-    if (ALGO == Algorithm::CN_CCX) {
+    if (props.isCCX()) {
         RESTORE_ROUNDING_MODE();
     }
 
@@ -1930,17 +1884,8 @@ inline void cryptonight_penta_hash(const uint8_t *__restrict__ input, size_t siz
 {
     constexpr CnAlgo<ALGO> props;
     constexpr size_t MASK        = props.mask();
-    constexpr Algorithm::Id BASE = props.base();
 
-#   ifdef XMRIG_ALGO_CN_HEAVY
-    constexpr bool IS_CN_HEAVY_TUBE = ALGO == Algorithm::CN_HEAVY_TUBE;
-    constexpr bool IS_CN_HEAVY_XHV  = ALGO == Algorithm::CN_HEAVY_XHV;
-#   else
-    constexpr bool IS_CN_HEAVY_TUBE = false;
-    constexpr bool IS_CN_HEAVY_XHV  = false;
-#   endif
-
-    if (BASE == Algorithm::CN_1 && size < 43) {
+    if (props.isBase1() && size < 43) {
         memset(output, 0, 32 * 5);
         return;
     }
@@ -1970,7 +1915,7 @@ inline void cryptonight_penta_hash(const uint8_t *__restrict__ input, size_t siz
     CONST_INIT(ctx[3], 3);
     CONST_INIT(ctx[4], 4);
     VARIANT2_SET_ROUNDING_MODE();
-    if (ALGO == Algorithm::CN_CCX) {
+    if (props.isCCX()) {
         RESTORE_ROUNDING_MODE();
     }