Merge branch 'dev'

xmrig · Jan 23, 2019 · c85c567 · c85c567
2 parents a10847c + ffef176
commit c85c567
Show file tree

Hide file tree

Showing 10 changed files with 138 additions and 19 deletions.
diff --git a/CHANGELOG.md b/CHANGELOG.md
@@ -1,3 +1,6 @@
+# v2.10.0
+- [#904](https://github.com/xmrig/xmrig/issues/904) Added new algorithm `cn-pico/trtl` (aliases `cryptonight-turtle`, `cn-trtl`) for upcoming TurtleCoin (TRTL) fork.
+
 # v2.9.4
 - [#913](https://github.com/xmrig/xmrig/issues/913) Fixed Masari (MSR) support (this update required for upcoming fork).
 

diff --git a/src/common/crypto/Algorithm.cpp b/src/common/crypto/Algorithm.cpp
@@ -6,7 +6,7 @@
  * Copyright 2016      Jay D Dee   <[email protected]>
  * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
  * Copyright 2018      Lee Clagett <https://github.com/vtnerd>
- * Copyright 2018      SChernykh   <https://github.com/SChernykh>
+ * Copyright 2018-2019 SChernykh   <https://github.com/SChernykh>
  * Copyright 2016-2019 XMRig       <https://github.com/xmrig>, <[email protected]>
  *
  *   This program is free software: you can redistribute it and/or modify
@@ -78,6 +78,14 @@ static AlgoData const algorithms[] = {
     { "cryptonight-heavy/xhv",  "cn-heavy/xhv",  xmrig::CRYPTONIGHT_HEAVY, xmrig::VARIANT_XHV  },
     { "cryptonight-heavy/tube", "cn-heavy/tube", xmrig::CRYPTONIGHT_HEAVY, xmrig::VARIANT_TUBE },
 #   endif
+
+#   ifndef XMRIG_NO_CN_PICO
+    { "cryptonight-pico/trtl",  "cn-pico/trtl",  xmrig::CRYPTONIGHT_PICO, xmrig::VARIANT_TRTL },
+    { "cryptonight-pico",       "cn-pico",       xmrig::CRYPTONIGHT_PICO, xmrig::VARIANT_TRTL },
+    { "cryptonight-turtle",     "cn-trtl",       xmrig::CRYPTONIGHT_PICO, xmrig::VARIANT_TRTL },
+    { "cryptonight-ultralite",  "cn-ultralite",  xmrig::CRYPTONIGHT_PICO, xmrig::VARIANT_TRTL },
+    { "cryptonight_turtle",     "cn_turtle",     xmrig::CRYPTONIGHT_PICO, xmrig::VARIANT_TRTL },
+#   endif
 };
 
 
@@ -111,7 +119,8 @@ static const char *variants[] = {
     "xao",
     "rto",
     "2",
-    "half"
+    "half",
+    "trtl"
 };
 
 
@@ -225,6 +234,10 @@ void xmrig::Algorithm::parseVariant(int variant)
 void xmrig::Algorithm::setAlgo(Algo algo)
 {
     m_algo = algo;
+
+    if (m_algo == CRYPTONIGHT_PICO && m_variant == VARIANT_AUTO) {
+        m_variant = xmrig::VARIANT_TRTL;
+    }
 }
 
 

diff --git a/src/common/xmrig.h b/src/common/xmrig.h
@@ -5,7 +5,7 @@
  * Copyright 2014-2016 Wolf9466    <https://github.com/OhGodAPet>
  * Copyright 2016      Jay D Dee   <[email protected]>
  * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
- * Copyright 2018      SChernykh   <https://github.com/SChernykh>
+ * Copyright 2018-2019 SChernykh   <https://github.com/SChernykh>
  * Copyright 2016-2019 XMRig       <https://github.com/xmrig>, <[email protected]>
  *
  *   This program is free software: you can redistribute it and/or modify
@@ -32,9 +32,11 @@ namespace xmrig
 
 enum Algo {
     INVALID_ALGO = -1,
-    CRYPTONIGHT,       /* CryptoNight (Monero) */
-    CRYPTONIGHT_LITE,  /* CryptoNight-Lite (AEON) */
-    CRYPTONIGHT_HEAVY  /* CryptoNight-Heavy (RYO) */
+    CRYPTONIGHT,        /* CryptoNight (2 MB) */
+    CRYPTONIGHT_LITE,   /* CryptoNight (1 MB) */
+    CRYPTONIGHT_HEAVY,  /* CryptoNight (4 MB) */
+    CRYPTONIGHT_PICO,   /* CryptoNight (256 KB) */
+    ALGO_MAX
 };
 
 
@@ -70,6 +72,7 @@ enum Variant {
     VARIANT_RTO  = 7,  // Modified CryptoNight variant 1 (Arto only)
     VARIANT_2    = 8,  // CryptoNight variant 2
     VARIANT_HALF = 9,  // CryptoNight variant 2 with half iterations (Masari/Stellite)
+    VARIANT_TRTL = 10, // CryptoNight Turtle (TRTL)
     VARIANT_MAX
 };
 

diff --git a/src/crypto/CryptoNight.cpp b/src/crypto/CryptoNight.cpp
@@ -67,7 +67,7 @@ CryptoNight::cn_hash_fun CryptoNight::fn(xmrig::Algo algorithm, xmrig::AlgoVerif
 
     assert(variant >= VARIANT_0 && variant < VARIANT_MAX);
 
-    static const cn_hash_fun func_table[VARIANT_MAX * 2 * 3] = {
+    static const cn_hash_fun func_table[] = {
         cryptonight_single_hash<CRYPTONIGHT, false, VARIANT_0>,
         cryptonight_single_hash<CRYPTONIGHT, true,  VARIANT_0>,
 
@@ -96,6 +96,8 @@ CryptoNight::cn_hash_fun CryptoNight::fn(xmrig::Algo algorithm, xmrig::AlgoVerif
         cryptonight_single_hash<CRYPTONIGHT, false, VARIANT_HALF>,
         cryptonight_single_hash<CRYPTONIGHT, true,  VARIANT_HALF>,
 
+        nullptr, nullptr, // VARIANT_TRTL
+
 #       ifndef XMRIG_NO_AEON
         cryptonight_single_hash<CRYPTONIGHT_LITE, false, VARIANT_0>,
         cryptonight_single_hash<CRYPTONIGHT_LITE, true,  VARIANT_0>,
@@ -111,12 +113,14 @@ CryptoNight::cn_hash_fun CryptoNight::fn(xmrig::Algo algorithm, xmrig::AlgoVerif
         nullptr, nullptr, // VARIANT_RTO
         nullptr, nullptr, // VARIANT_2
         nullptr, nullptr, // VARIANT_HALF
+        nullptr, nullptr, // VARIANT_TRTL
 #       else
         nullptr, nullptr, nullptr, nullptr,
         nullptr, nullptr, nullptr, nullptr,
         nullptr, nullptr, nullptr, nullptr,
         nullptr, nullptr, nullptr, nullptr,
         nullptr, nullptr, nullptr, nullptr,
+        nullptr, nullptr,
 #       endif
 
 #       ifndef XMRIG_NO_SUMO
@@ -138,15 +142,41 @@ CryptoNight::cn_hash_fun CryptoNight::fn(xmrig::Algo algorithm, xmrig::AlgoVerif
         nullptr, nullptr, // VARIANT_RTO
         nullptr, nullptr, // VARIANT_2
         nullptr, nullptr, // VARIANT_HALF
+        nullptr, nullptr, // VARIANT_TRTL
 #       else
         nullptr, nullptr, nullptr, nullptr,
         nullptr, nullptr, nullptr, nullptr,
         nullptr, nullptr, nullptr, nullptr,
         nullptr, nullptr, nullptr, nullptr,
         nullptr, nullptr, nullptr, nullptr,
+        nullptr, nullptr,
+#       endif
+#       ifndef XMRIG_NO_CN_PICO
+        nullptr, nullptr, // VARIANT_0
+        nullptr, nullptr, // VARIANT_1
+        nullptr, nullptr, // VARIANT_TUBE
+        nullptr, nullptr, // VARIANT_XTL
+        nullptr, nullptr, // VARIANT_MSR
+        nullptr, nullptr, // VARIANT_XHV
+        nullptr, nullptr, // VARIANT_XAO
+        nullptr, nullptr, // VARIANT_RTO
+        nullptr, nullptr, // VARIANT_2
+        nullptr, nullptr, // VARIANT_HALF
+
+        cryptonight_single_hash<CRYPTONIGHT_PICO, false, VARIANT_TRTL>,
+        cryptonight_single_hash<CRYPTONIGHT_PICO, true,  VARIANT_TRTL>,
+    #else
+        nullptr, nullptr, nullptr, nullptr,
+        nullptr, nullptr, nullptr, nullptr,
+        nullptr, nullptr, nullptr, nullptr,
+        nullptr, nullptr, nullptr, nullptr,
+        nullptr, nullptr, nullptr, nullptr,
+        nullptr, nullptr,
 #       endif
     };
 
+    static_assert((VARIANT_MAX * 2 * ALGO_MAX) == sizeof(func_table) / sizeof(func_table[0]), "func_table size mismatch");
+
     const size_t index = VARIANT_MAX * 2 * algorithm + 2 * variant + av - 1;
 
 #   ifndef NDEBUG
@@ -209,6 +239,12 @@ bool CryptoNight::selfTest() {
     }
 #   endif
 
+#   ifndef XMRIG_NO_CN_PICO
+    if (m_algorithm == xmrig::CRYPTONIGHT_PICO) {
+        return verify(VARIANT_TRTL, test_output_pico_trtl);
+    }
+#   endif
+
     return false;
 }
 

diff --git a/src/crypto/CryptoNight_constants.h b/src/crypto/CryptoNight_constants.h
@@ -6,6 +6,7 @@
  * Copyright 2016      Jay D Dee   <[email protected]>
  * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
  * Copyright 2018      Lee Clagett <https://github.com/vtnerd>
+ * Copyright 2018-2019 SChernykh   <https://github.com/SChernykh>
  * Copyright 2016-2019 XMRig       <https://github.com/xmrig>, <[email protected]>
  *
  *   This program is free software: you can redistribute it and/or modify
@@ -22,8 +23,8 @@
  *   along with this program. If not, see <http://www.gnu.org/licenses/>.
  */
 
-#ifndef __CRYPTONIGHT_CONSTANTS_H__
-#define __CRYPTONIGHT_CONSTANTS_H__
+#ifndef XMRIG_CRYPTONIGHT_CONSTANTS_H
+#define XMRIG_CRYPTONIGHT_CONSTANTS_H
 
 
 #include <stdint.h>
@@ -49,11 +50,17 @@ constexpr const size_t   CRYPTONIGHT_HEAVY_MEMORY = 4 * 1024 * 1024;
 constexpr const uint32_t CRYPTONIGHT_HEAVY_MASK   = 0x3FFFF0;
 constexpr const uint32_t CRYPTONIGHT_HEAVY_ITER   = 0x40000;
 
+constexpr const size_t   CRYPTONIGHT_PICO_MEMORY = 256 * 1024;
+constexpr const uint32_t CRYPTONIGHT_PICO_MASK   = 0x1FFF0;
+constexpr const uint32_t CRYPTONIGHT_PICO_ITER   = 0x40000;
+constexpr const uint32_t CRYPTONIGHT_TRTL_ITER   = 0x10000;
+
 
 template<Algo ALGO> inline constexpr size_t cn_select_memory()           { return 0; }
 template<> inline constexpr size_t cn_select_memory<CRYPTONIGHT>()       { return CRYPTONIGHT_MEMORY; }
 template<> inline constexpr size_t cn_select_memory<CRYPTONIGHT_LITE>()  { return CRYPTONIGHT_LITE_MEMORY; }
 template<> inline constexpr size_t cn_select_memory<CRYPTONIGHT_HEAVY>() { return CRYPTONIGHT_HEAVY_MEMORY; }
+template<> inline constexpr size_t cn_select_memory<CRYPTONIGHT_PICO>()  { return CRYPTONIGHT_PICO_MEMORY; }
 
 
 inline size_t cn_select_memory(Algo algorithm)
@@ -69,6 +76,9 @@ inline size_t cn_select_memory(Algo algorithm)
     case CRYPTONIGHT_HEAVY:
         return CRYPTONIGHT_HEAVY_MEMORY;
 
+    case CRYPTONIGHT_PICO:
+        return CRYPTONIGHT_PICO_MEMORY;
+
     default:
         break;
     }
@@ -81,6 +91,7 @@ template<Algo ALGO> inline constexpr uint32_t cn_select_mask()           { retur
 template<> inline constexpr uint32_t cn_select_mask<CRYPTONIGHT>()       { return CRYPTONIGHT_MASK; }
 template<> inline constexpr uint32_t cn_select_mask<CRYPTONIGHT_LITE>()  { return CRYPTONIGHT_LITE_MASK; }
 template<> inline constexpr uint32_t cn_select_mask<CRYPTONIGHT_HEAVY>() { return CRYPTONIGHT_HEAVY_MASK; }
+template<> inline constexpr uint32_t cn_select_mask<CRYPTONIGHT_PICO>()  { return CRYPTONIGHT_PICO_MASK; }
 
 
 inline uint32_t cn_select_mask(Algo algorithm)
@@ -96,6 +107,9 @@ inline uint32_t cn_select_mask(Algo algorithm)
     case CRYPTONIGHT_HEAVY:
         return CRYPTONIGHT_HEAVY_MASK;
 
+    case CRYPTONIGHT_PICO:
+        return CRYPTONIGHT_PICO_MASK;
+
     default:
         break;
     }
@@ -118,6 +132,7 @@ template<> inline constexpr uint32_t cn_select_iter<CRYPTONIGHT_LITE, VARIANT_1>
 template<> inline constexpr uint32_t cn_select_iter<CRYPTONIGHT_HEAVY, VARIANT_0>()    { return CRYPTONIGHT_HEAVY_ITER; }
 template<> inline constexpr uint32_t cn_select_iter<CRYPTONIGHT_HEAVY, VARIANT_XHV>()  { return CRYPTONIGHT_HEAVY_ITER; }
 template<> inline constexpr uint32_t cn_select_iter<CRYPTONIGHT_HEAVY, VARIANT_TUBE>() { return CRYPTONIGHT_HEAVY_ITER; }
+template<> inline constexpr uint32_t cn_select_iter<CRYPTONIGHT_PICO, VARIANT_TRTL>()  { return CRYPTONIGHT_TRTL_ITER; }
 
 
 inline uint32_t cn_select_iter(Algo algorithm, Variant variant)
@@ -130,6 +145,9 @@ inline uint32_t cn_select_iter(Algo algorithm, Variant variant)
     case VARIANT_RTO:
         return CRYPTONIGHT_XAO_ITER;
 
+    case VARIANT_TRTL:
+        return CRYPTONIGHT_TRTL_ITER;
+
     default:
         break;
     }
@@ -145,6 +163,9 @@ inline uint32_t cn_select_iter(Algo algorithm, Variant variant)
     case CRYPTONIGHT_HEAVY:
         return CRYPTONIGHT_HEAVY_ITER;
 
+    case CRYPTONIGHT_PICO:
+        return CRYPTONIGHT_TRTL_ITER;
+
     default:
         break;
     }
@@ -164,9 +185,10 @@ template<> inline constexpr Variant cn_base_variant<VARIANT_XAO>()   { return VA
 template<> inline constexpr Variant cn_base_variant<VARIANT_RTO>()   { return VARIANT_1; }
 template<> inline constexpr Variant cn_base_variant<VARIANT_2>()     { return VARIANT_2; }
 template<> inline constexpr Variant cn_base_variant<VARIANT_HALF>()  { return VARIANT_2; }
+template<> inline constexpr Variant cn_base_variant<VARIANT_TRTL>()  { return VARIANT_2; }
 
 
 } /* namespace xmrig */
 
 
-#endif /* __CRYPTONIGHT_CONSTANTS_H__ */
+#endif /* XMRIG_CRYPTONIGHT_CONSTANTS_H */
diff --git a/src/crypto/CryptoNight_test.h b/src/crypto/CryptoNight_test.h
@@ -6,6 +6,7 @@
  * Copyright 2016      Jay D Dee   <[email protected]>
  * Copyright 2017-2018 XMR-Stak    <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
  * Copyright 2018      Lee Clagett <https://github.com/vtnerd>
+ * Copyright 2018-2019 SChernykh   <https://github.com/SChernykh>
  * Copyright 2016-2019 XMRig       <https://github.com/xmrig>, <[email protected]>
  *
  *   This program is free software: you can redistribute it and/or modify
@@ -254,4 +255,21 @@ const static uint8_t test_output_tube_heavy[160] = {
 #endif
 
 
+#ifndef XMRIG_NO_CN_PICO
+// "cn-pico/trtl"
+const static uint8_t test_output_pico_trtl[160] = {
+    0x08, 0xF4, 0x21, 0xD7, 0x83, 0x31, 0x17, 0x30, 0x0E, 0xDA, 0x66, 0xE9, 0x8F, 0x4A, 0x25, 0x69,
+    0x09, 0x3D, 0xF3, 0x00, 0x50, 0x01, 0x73, 0x94, 0x4E, 0xFC, 0x40, 0x1E, 0x9A, 0x4A, 0x17, 0xAF,
+    0xB2, 0x17, 0x2E, 0xC9, 0x46, 0x6E, 0x1A, 0xEE, 0x70, 0xEC, 0x85, 0x72, 0xA1, 0x4C, 0x23, 0x3E,
+    0xE3, 0x54, 0x58, 0x2B, 0xCB, 0x93, 0xF8, 0x69, 0xD4, 0x29, 0x74, 0x4D, 0xE5, 0x72, 0x6A, 0x26,
+    0x4E, 0xFD, 0x28, 0xFC, 0xD3, 0x74, 0x8A, 0x83, 0xF3, 0xCA, 0x92, 0x84, 0xE7, 0x4E, 0x10, 0xC2,
+    0x05, 0x62, 0xC7, 0xBE, 0x99, 0x73, 0xED, 0x90, 0xB5, 0x6F, 0xDA, 0x64, 0x71, 0x2D, 0x99, 0x39,
+    0x29, 0xDB, 0x22, 0x2B, 0x97, 0xB6, 0x37, 0x0E, 0x9A, 0x03, 0x65, 0xCC, 0xF7, 0xD0, 0x9A, 0xB7,
+    0x68, 0xCE, 0x07, 0x3E, 0x15, 0x40, 0x3C, 0xCE, 0x8C, 0x63, 0x16, 0x72, 0xB5, 0x74, 0x84, 0xF4,
+    0xA1, 0xE7, 0x53, 0x85, 0xFB, 0x72, 0xDD, 0x75, 0x90, 0x39, 0xB2, 0x3D, 0xC3, 0x08, 0x2C, 0xD5,
+    0x01, 0x08, 0x27, 0x75, 0x86, 0xB9, 0xBB, 0x9B, 0xDF, 0xEA, 0x49, 0xDE, 0x46, 0xCB, 0x83, 0x45
+};
+#endif
+
+
 #endif /* XMRIG_CRYPTONIGHT_TEST_H */
diff --git a/src/crypto/CryptoNight_x86.h b/src/crypto/CryptoNight_x86.h
@@ -458,6 +458,8 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si
     constexpr size_t MEM          = xmrig::cn_select_memory<ALGO>();
     constexpr xmrig::Variant BASE = xmrig::cn_base_variant<VARIANT>();
 
+    static_assert(MASK > 0 && ITERATIONS > 0 && MEM > 0, "unsupported algorithm/variant");
+
     if (BASE == xmrig::VARIANT_1 && size < 43) {
         memset(output, 0, 32);
         return;
@@ -570,11 +572,16 @@ extern "C" void cnv2_mainloop_ryzen_asm(cryptonight_ctx *ctx);
 extern "C" void cnv2_mainloop_bulldozer_asm(cryptonight_ctx *ctx);
 extern "C" void cnv2_double_mainloop_sandybridge_asm(cryptonight_ctx *ctx0, cryptonight_ctx *ctx1);
 
-extern xmrig::CpuThread::cn_mainloop_fun cn_half_mainloop_ivybridge_asm;
-extern xmrig::CpuThread::cn_mainloop_fun cn_half_mainloop_ryzen_asm;
-extern xmrig::CpuThread::cn_mainloop_fun cn_half_mainloop_bulldozer_asm;
+extern xmrig::CpuThread::cn_mainloop_fun        cn_half_mainloop_ivybridge_asm;
+extern xmrig::CpuThread::cn_mainloop_fun        cn_half_mainloop_ryzen_asm;
+extern xmrig::CpuThread::cn_mainloop_fun        cn_half_mainloop_bulldozer_asm;
 extern xmrig::CpuThread::cn_mainloop_double_fun cn_half_double_mainloop_sandybridge_asm;
 
+extern xmrig::CpuThread::cn_mainloop_fun        cn_trtl_mainloop_ivybridge_asm;
+extern xmrig::CpuThread::cn_mainloop_fun        cn_trtl_mainloop_ryzen_asm;
+extern xmrig::CpuThread::cn_mainloop_fun        cn_trtl_mainloop_bulldozer_asm;
+extern xmrig::CpuThread::cn_mainloop_double_fun cn_trtl_double_mainloop_sandybridge_asm;
+
 
 template<xmrig::Algo ALGO, xmrig::Variant VARIANT, xmrig::Assembly ASM>
 inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_t size, uint8_t *__restrict__ output, cryptonight_ctx **__restrict__ ctx)
@@ -606,6 +613,17 @@ inline void cryptonight_single_hash_asm(const uint8_t *__restrict__ input, size_
             cn_half_mainloop_bulldozer_asm(ctx[0]);
         }
     }
+    else if (VARIANT == xmrig::VARIANT_TRTL) {
+        if (ASM == xmrig::ASM_INTEL) {
+            cn_trtl_mainloop_ivybridge_asm(ctx[0]);
+        }
+        else if (ASM == xmrig::ASM_RYZEN) {
+            cn_trtl_mainloop_ryzen_asm(ctx[0]);
+        }
+        else {
+            cn_trtl_mainloop_bulldozer_asm(ctx[0]);
+        }
+    }
 
     cn_implode_scratchpad<ALGO, MEM, false>(reinterpret_cast<__m128i*>(ctx[0]->memory), reinterpret_cast<__m128i*>(ctx[0]->state));
     xmrig::keccakf(reinterpret_cast<uint64_t*>(ctx[0]->state), 24);
@@ -630,6 +648,9 @@ inline void cryptonight_double_hash_asm(const uint8_t *__restrict__ input, size_
     else if (VARIANT == xmrig::VARIANT_HALF) {
         cn_half_double_mainloop_sandybridge_asm(ctx[0], ctx[1]);
     }
+    else if (VARIANT == xmrig::VARIANT_TRTL) {
+        cn_trtl_double_mainloop_sandybridge_asm(ctx[0], ctx[1]);
+    }
 
     cn_implode_scratchpad<ALGO, MEM, false>(reinterpret_cast<__m128i*>(ctx[0]->memory), reinterpret_cast<__m128i*>(ctx[0]->state));
     cn_implode_scratchpad<ALGO, MEM, false>(reinterpret_cast<__m128i*>(ctx[1]->memory), reinterpret_cast<__m128i*>(ctx[1]->state));

diff --git a/src/nvidia/cuda_core.cu b/src/nvidia/cuda_core.cu
@@ -329,7 +329,7 @@ __global__ void cryptonight_core_gpu_phase2_double(
 
     uint64_t* ptr0;
     for (int i = start; i < end; ++i) {
-        ptr0 = (uint64_t *)&l0[idx0 & 0x1FFFC0];
+        ptr0 = (uint64_t *)&l0[idx0 & (MASK - 0x30)];
 
         ((ulonglong4*)myChunks)[sub] = ((ulonglong4*)ptr0)[sub];
 
@@ -364,7 +364,7 @@ __global__ void cryptonight_core_gpu_phase2_double(
 
         idx0 = shuffle<2>(sPtr, sub, cx_aes.x, 0);
         idx1 = (idx0 & 0x30) >> 3;
-        ptr0 = (uint64_t *)&l0[idx0 & MASK & 0x1FFFC0];
+        ptr0 = (uint64_t *)&l0[idx0 & MASK & (MASK - 0x30)];
 
         ((ulonglong4*)myChunks)[sub] = ((ulonglong4*)ptr0)[sub];
 
@@ -852,4 +852,7 @@ void cryptonight_gpu_hash(nvid_ctx *ctx, xmrig::Algo algo, xmrig::Variant varian
             break;
         }
     }
+    else if (algo == CRYPTONIGHT_PICO) {
+        cryptonight_core_gpu_hash<CRYPTONIGHT_PICO, VARIANT_TRTL>(ctx, startNonce);
+    }
 }