Fix mismatch in CN_GPU algorithm id (upstream changed to 0300 vs 00ff)

MoneroOcean · Dec 17, 2021 · 5836da6 · 5836da6
1 parent 0a5ebf5
commit 5836da6
Show file tree

Hide file tree

Showing 6 changed files with 492 additions and 472 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -10,8 +10,8 @@ option(WITH_CN_R       "Enable CryptoNight-R algorithm" ON)
 option(WITH_CN_LITE    "Enable CryptoNight-Lite algorithms family" ON)
 option(WITH_CN_HEAVY   "Enable CryptoNight-Heavy algorithms family" ON)
 option(WITH_CN_PICO    "Enable CryptoNight-Pico algorithm" ON)
-option(WITH_CN_GPU     "Enable CryptoNight-GPU algorithm" ON)
 option(WITH_CN_FEMTO   "Enable CryptoNight-UPX2 algorithm" ON)
+option(WITH_CN_GPU     "Enable CryptoNight-GPU algorithm" ON)
 option(WITH_ARGON2     "Enable Argon2 algorithms family" OFF) #unsupported
 
 if (CUDA_VERSION VERSION_LESS 9.0)
@@ -37,14 +37,14 @@ if (WITH_CN_PICO)
     add_definitions(/DXMRIG_ALGO_CN_PICO)
 endif()
 
-if (WITH_CN_GPU)
-    add_definitions(/DXMRIG_ALGO_CN_GPU)
-endif()
-
 if (WITH_CN_FEMTO)
     add_definitions(/DXMRIG_ALGO_CN_FEMTO)
 endif()
 
+if (WITH_CN_GPU)
+    add_definitions(/DXMRIG_ALGO_CN_GPU)
+endif()
+
 if (WITH_RANDOMX)
     add_definitions(/DXMRIG_ALGO_RANDOMX)
 endif()

diff --git a/src/crypto/cn/CnAlgo.h b/src/crypto/cn/CnAlgo.h
@@ -82,16 +82,16 @@ class CnAlgo
             return CN_ITER / 8;
 #       endif
 
-#       ifdef XMRIG_ALGO_CN_GPU
-        case Algorithm::CN_GPU:
-            return 0xC000;
-#	endif
-
 #       ifdef XMRIG_ALGO_CN_FEMTO
         case Algorithm::CN_UPX2:
             return CN_ITER / 32;
 #       endif
 
+#       ifdef XMRIG_ALGO_CN_GPU
+        case Algorithm::CN_GPU:
+            return 0xC000;
+#       endif
+
         default:
             break;
         }
@@ -107,18 +107,18 @@ class CnAlgo
         }
 #       endif
 
-#       ifdef XMRIG_ALGO_CN_GPU
-        if (algo == Algorithm::CN_GPU) {
-            return 0x1FFFC0;
-        }
-#       endif
-
 #       ifdef XMRIG_ALGO_CN_FEMTO
         if (algo == Algorithm::CN_UPX2) {
             return 0x1FFF0;
         }
 #       endif
 
+#       ifdef XMRIG_ALGO_CN_GPU
+        if (algo == Algorithm::CN_GPU) {
+            return 0x1FFFC0;
+        }
+#       endif
+
         return ((Algorithm::l3(algo) - 1) / 16) * 16;
     }
 
@@ -140,14 +140,14 @@ template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_RWZ>::iterations() con
 template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_ZLS>::iterations() const          { return 0x60000; }
 template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_PICO_0>::iterations() const       { return CN_ITER / 8; }
 template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_PICO_TLO>::iterations() const     { return CN_ITER / 8; }
-template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_GPU>::iterations() const          { return 0xC000; }
 template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_CCX>::iterations() const          { return CN_ITER / 2; }
 template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_UPX2>::iterations() const         { return CN_ITER / 32; }
+template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_GPU>::iterations() const          { return 0xC000; }
 
 
 template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_PICO_0>::mask() const             { return 0x1FFF0; }
-template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_GPU>::mask() const                { return 0x1FFFC0; }
 template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_UPX2>::mask() const               { return 0x1FFF0; }
+template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_GPU>::mask() const                { return 0x1FFFC0; }
 
 
 } /* namespace xmrig_cuda */

diff --git a/src/crypto/common/Algorithm.h b/src/crypto/common/Algorithm.h
@@ -21,6 +21,7 @@
 #define XMRIG_ALGORITHM_H
 
 
+#include <cstddef>
 #include <cstdint>
 #include <vector>
 
@@ -53,6 +54,7 @@ class Algorithm
         CN_PICO_0       = 0x63120200,   // "cn-pico"          CryptoNight-Pico
         CN_PICO_TLO     = 0x63120274,   // "cn-pico/tlo"      CryptoNight-Pico (TLO)
         CN_UPX2         = 0x63110200,   // "cn/upx2"          Uplexa (UPX2)
+        CN_GPU          = 0x63150300,   // "cn/gpu"           CryptoNight-GPU (Ryo).
         RX_0            = 0x72151200,   // "rx/0"             RandomX (reference configuration).
         RX_WOW          = 0x72141177,   // "rx/wow"           RandomWOW (Wownero).
         RX_ARQ          = 0x72121061,   // "rx/arq"           RandomARQ (Arqma).
@@ -65,7 +67,6 @@ class Algorithm
         ASTROBWT_DERO   = 0x41000000,   // "astrobwt"         AstroBWT (Dero)
         KAWPOW_RVN      = 0x6b0f0000,   // "kawpow/rvn"       KawPow (RVN)
 
-        CN_GPU          = 0x631500ff,   // "cn/gpu"           CryptoNight-GPU (Ryo).
         RX_XLA          = 0x721211ff,   // "panthera"         Panthera (Scala2).
     };
 

diff --git a/src/cuda_core.cu b/src/cuda_core.cu
@@ -837,7 +837,7 @@ void cryptonight_core_gpu_hash_gpu(nvid_ctx* ctx, uint32_t nonce)
 
     CUDA_CHECK_KERNEL(
         ctx->device_id,
-        xmrig::cn_gpu::cn_explode_gpu<MEMORY><<<intensity,32>>>((int*)ctx->d_ctx_state, (int*)ctx->d_long_state)
+        xmrig_cuda::cn_gpu::cn_explode_gpu<MEMORY><<<intensity,32>>>((int*)ctx->d_ctx_state, (int*)ctx->d_long_state)
     );
 
     int partcount = 1 << ctx->device_bfactor;
@@ -846,7 +846,7 @@ void cryptonight_core_gpu_hash_gpu(nvid_ctx* ctx, uint32_t nonce)
         CUDA_CHECK_KERNEL(
             ctx->device_id,
             // 36 x 16byte x numThreads
-            xmrig::cn_gpu::cryptonight_core_gpu_phase2_gpu<ITERATIONS, MEMORY>
+            xmrig_cuda::cn_gpu::cryptonight_core_gpu_phase2_gpu<ITERATIONS, MEMORY>
                 <<<ctx->device_blocks, ctx->device_threads * 16,  36 * 16 * ctx->device_threads>>>
                 (
                     (int*)ctx->d_ctx_state,