diff --git a/CMakeLists.txt b/CMakeLists.txt index 12712fe1c8..0e6ed895c1 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -26,7 +26,6 @@ option(WITH_NVML "Enable NVML (NVIDIA Management Library) support (on option(WITH_ADL "Enable ADL (AMD Display Library) or sysfs support (only if OpenCL backend enabled)" ON) option(WITH_STRICT_CACHE "Enable strict checks for OpenCL cache" ON) option(WITH_INTERLEAVE_DEBUG_LOG "Enable debug log for threads interleave" OFF) -option(WITH_MO_BENCHMARK "Enable Benchmark module and algo-perf feature (for MoneroOcean)" ON) option(WITH_PROFILING "Enable profiling for developers" OFF) option(WITH_SSE4_1 "Enable SSE 4.1 for Blake2" ON) option(WITH_AVX2 "Enable AVX2 for Blake2" ON) @@ -34,6 +33,7 @@ option(WITH_VAES "Enable VAES instructions for Cryptonight" ON) option(WITH_BENCHMARK "Enable builtin RandomX benchmark and stress test" ON) option(WITH_SECURE_JIT "Enable secure access to JIT memory" OFF) option(WITH_DMI "Enable DMI/SMBIOS reader" ON) +option(WITH_MO_BENCHMARK "Enable Benchmark module and algo-perf feature (for MoneroOcean)" ON) option(BUILD_STATIC "Build static binary" OFF) option(ARM_V8 "Force ARMv8 (64 bit) architecture, use with caution if automatic detection fails, but you sure it may work" OFF) @@ -135,13 +135,6 @@ set(SOURCES_CRYPTO src/crypto/flex/flex_keccak.c ) -if (WITH_MO_BENCHMARK) - list(APPEND SOURCES - src/core/MoBenchmark.cpp - ) - add_definitions(/DXMRIG_FEATURE_MO_BENCHMARK) -endif() - if (CMAKE_C_COMPILER_ID MATCHES GNU) set_source_files_properties(src/crypto/cn/CnHash.cpp PROPERTIES COMPILE_FLAGS "-Ofast -fno-tree-vectorize") endif() @@ -248,6 +241,13 @@ if (WITH_DEBUG_LOG) add_definitions(/DAPP_DEBUG) endif() +if (WITH_MO_BENCHMARK) + list(APPEND SOURCES + src/core/MoBenchmark.cpp + ) + add_definitions(/DXMRIG_FEATURE_MO_BENCHMARK) +endif() + add_executable(${CMAKE_PROJECT_NAME} ${HEADERS} ${SOURCES} ${SOURCES_OS} ${HEADERS_CRYPTO} ${SOURCES_CRYPTO} ${SOURCES_SYSLOG} ${TLS_SOURCES} ${XMRIG_ASM_SOURCES}) target_link_libraries(${CMAKE_PROJECT_NAME} ${XMRIG_ASM_LIBRARY} ${OPENSSL_LIBRARIES} ${UV_LIBRARIES} ${EXTRA_LIBS} ${CPUID_LIB} ${ARGON2_LIBRARY} ${ETHASH_LIBRARY} ${GHOSTRIDER_LIBRARY}) diff --git a/cmake/randomx.cmake b/cmake/randomx.cmake index ec98f340ba..f24ed50502 100644 --- a/cmake/randomx.cmake +++ b/cmake/randomx.cmake @@ -55,12 +55,12 @@ if (WITH_RANDOMX) src/crypto/rx/RxQueue.cpp src/crypto/rx/RxVm.cpp - ### Removed useless includes - src/crypto/randomx/panthera/sha256.c - src/crypto/randomx/panthera/KangarooTwelve.c - src/crypto/randomx/panthera/KeccakP-1600-reference.c - src/crypto/randomx/panthera/KeccakSpongeWidth1600.c - src/crypto/randomx/panthera/yespower-opt.c + ### Removed useless includes + src/crypto/randomx/panthera/sha256.c + src/crypto/randomx/panthera/KangarooTwelve.c + src/crypto/randomx/panthera/KeccakP-1600-reference.c + src/crypto/randomx/panthera/KeccakSpongeWidth1600.c + src/crypto/randomx/panthera/yespower-opt.c ) if (WITH_ASM AND CMAKE_C_COMPILER_ID MATCHES MSVC) diff --git a/src/backend/common/Tags.h b/src/backend/common/Tags.h index b77891791a..0e29827595 100644 --- a/src/backend/common/Tags.h +++ b/src/backend/common/Tags.h @@ -46,10 +46,6 @@ const char *ocl_tag(); const char *cuda_tag(); #endif -#ifdef XMRIG_FEATURE_MO_BENCHMARK -const char *bm_tag(); -#endif - } // namespace xmrig diff --git a/src/backend/common/Threads.cpp b/src/backend/common/Threads.cpp index 4b826fe9b5..2931afeb97 100644 --- a/src/backend/common/Threads.cpp +++ b/src/backend/common/Threads.cpp @@ -128,7 +128,9 @@ xmrig::String xmrig::Threads::profileName(const Algorithm &algorithm, bool st } } - if (std::is_same::value && (name == "defyx" || name == "panthera") && has("rx")) return "rx"; + if (std::is_same::value && (name == Algorithm::kRX_XLA) && has(Algorithm::kRX)) { + return Algorithm::kRX; + } if (has(kAsterisk)) { return kAsterisk; diff --git a/src/backend/cpu/platform/BasicCpuInfo.cpp b/src/backend/cpu/platform/BasicCpuInfo.cpp index c8a7205a13..408b355222 100644 --- a/src/backend/cpu/platform/BasicCpuInfo.cpp +++ b/src/backend/cpu/platform/BasicCpuInfo.cpp @@ -350,7 +350,7 @@ xmrig::CpuThreads xmrig::BasicCpuInfo::threads(const Algorithm &algorithm, uint3 # ifdef XMRIG_ALGO_CN_FEMTO if (f == Algorithm::CN_FEMTO) { - return CpuThreads(count, 2); + return CpuThreads(count_limit, 2); } # endif diff --git a/src/backend/opencl/cl/cn/algorithm.cl b/src/backend/opencl/cl/cn/algorithm.cl index 25651a8908..8c091c46c9 100644 --- a/src/backend/opencl/cl/cn/algorithm.cl +++ b/src/backend/opencl/cl/cn/algorithm.cl @@ -24,6 +24,7 @@ #define ALGO_RX_WOW 0x72141177 #define ALGO_RX_ARQ 0x72121061 #define ALGO_RX_XEQ 0x72121000 +#define ALGO_RX_XLA 0x721211ff #define ALGO_RX_SFX 0x72151273 #define ALGO_RX_GRAFT 0x72151267 #define ALGO_RX_YADA 0x72151279 @@ -32,8 +33,6 @@ #define ALGO_AR2_WRKZ 0x61120000 #define ALGO_KAWPOW_RVN 0x6b0f0000 -#define ALGO_RX_XLA 0x721211ff - #define FAMILY_UNKNOWN 0 #define FAMILY_CN 0x63150000 #define FAMILY_CN_LITE 0x63140000 diff --git a/src/backend/opencl/cl/cn/cryptonight_cl.h b/src/backend/opencl/cl/cn/cryptonight_cl.h index cfc44e73f2..87942f31f0 100644 --- a/src/backend/opencl/cl/cn/cryptonight_cl.h +++ b/src/backend/opencl/cl/cn/cryptonight_cl.h @@ -36,15 +36,15 @@ static const char cryptonight_cl[61568] = { 0x32,0x30,0x32,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x57,0x4f,0x57,0x20,0x30,0x78,0x37,0x32,0x31,0x34,0x31,0x31, 0x37,0x37,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x41,0x52,0x51,0x20,0x30,0x78,0x37,0x32,0x31,0x32,0x31,0x30,0x36, 0x31,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x58,0x45,0x51,0x20,0x30,0x78,0x37,0x32,0x31,0x32,0x31,0x30,0x30,0x30, - 0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x53,0x46,0x58,0x20,0x30,0x78,0x37,0x32,0x31,0x35,0x31,0x32,0x37,0x33,0x0a, - 0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x47,0x52,0x41,0x46,0x54,0x20,0x30,0x78,0x37,0x32,0x31,0x35,0x31,0x32,0x36,0x37, - 0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x59,0x41,0x44,0x41,0x20,0x30,0x78,0x37,0x32,0x31,0x35,0x31,0x32,0x37,0x39, - 0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,0x41,0x20,0x30,0x78,0x36,0x31,0x31,0x33,0x30, - 0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,0x41,0x5f,0x56,0x32,0x20,0x30, - 0x78,0x36,0x31,0x31,0x34,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x57,0x52,0x4b,0x5a,0x20, - 0x30,0x78,0x36,0x31,0x31,0x32,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x4b,0x41,0x57,0x50,0x4f,0x57,0x5f,0x52, - 0x56,0x4e,0x20,0x30,0x78,0x36,0x62,0x30,0x66,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x58,0x4c, - 0x41,0x20,0x30,0x78,0x37,0x32,0x31,0x32,0x31,0x31,0x66,0x66,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x41,0x4d,0x49,0x4c,0x59,0x5f,0x55,0x4e,0x4b,0x4e, + 0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x58,0x4c,0x41,0x20,0x30,0x78,0x37,0x32,0x31,0x32,0x31,0x31,0x66,0x66,0x0a, + 0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x53,0x46,0x58,0x20,0x30,0x78,0x37,0x32,0x31,0x35,0x31,0x32,0x37,0x33,0x0a,0x23, + 0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x47,0x52,0x41,0x46,0x54,0x20,0x30,0x78,0x37,0x32,0x31,0x35,0x31,0x32,0x36,0x37,0x0a, + 0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x59,0x41,0x44,0x41,0x20,0x30,0x78,0x37,0x32,0x31,0x35,0x31,0x32,0x37,0x39,0x0a, + 0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,0x41,0x20,0x30,0x78,0x36,0x31,0x31,0x33,0x30,0x30, + 0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,0x41,0x5f,0x56,0x32,0x20,0x30,0x78, + 0x36,0x31,0x31,0x34,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x57,0x52,0x4b,0x5a,0x20,0x30, + 0x78,0x36,0x31,0x31,0x32,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x4b,0x41,0x57,0x50,0x4f,0x57,0x5f,0x52,0x56, + 0x4e,0x20,0x30,0x78,0x36,0x62,0x30,0x66,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x41,0x4d,0x49,0x4c,0x59,0x5f,0x55,0x4e,0x4b,0x4e, 0x4f,0x57,0x4e,0x20,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x41,0x4d,0x49,0x4c,0x59,0x5f,0x43,0x4e,0x20,0x30,0x78,0x36,0x33,0x31,0x35,0x30,0x30, 0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x41,0x4d,0x49,0x4c,0x59,0x5f,0x43,0x4e,0x5f,0x4c,0x49,0x54,0x45,0x20,0x30,0x78,0x36,0x33,0x31,0x34, 0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x41,0x4d,0x49,0x4c,0x59,0x5f,0x43,0x4e,0x5f,0x48,0x45,0x41,0x56,0x59,0x20,0x30,0x78,0x36, diff --git a/src/backend/opencl/cl/rx/randomx_cl.h b/src/backend/opencl/cl/rx/randomx_cl.h index bd160f3868..64a2cb884f 100644 --- a/src/backend/opencl/cl/rx/randomx_cl.h +++ b/src/backend/opencl/cl/rx/randomx_cl.h @@ -29,15 +29,15 @@ static const char randomx_cl[137870] = { 0x30,0x32,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x57,0x4f,0x57,0x20,0x30,0x78,0x37,0x32,0x31,0x34,0x31,0x31,0x37, 0x37,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x41,0x52,0x51,0x20,0x30,0x78,0x37,0x32,0x31,0x32,0x31,0x30,0x36,0x31, 0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x58,0x45,0x51,0x20,0x30,0x78,0x37,0x32,0x31,0x32,0x31,0x30,0x30,0x30,0x0a, - 0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x53,0x46,0x58,0x20,0x30,0x78,0x37,0x32,0x31,0x35,0x31,0x32,0x37,0x33,0x0a,0x23, - 0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x47,0x52,0x41,0x46,0x54,0x20,0x30,0x78,0x37,0x32,0x31,0x35,0x31,0x32,0x36,0x37,0x0a, - 0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x59,0x41,0x44,0x41,0x20,0x30,0x78,0x37,0x32,0x31,0x35,0x31,0x32,0x37,0x39,0x0a, - 0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,0x41,0x20,0x30,0x78,0x36,0x31,0x31,0x33,0x30,0x30, - 0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,0x41,0x5f,0x56,0x32,0x20,0x30,0x78, - 0x36,0x31,0x31,0x34,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x57,0x52,0x4b,0x5a,0x20,0x30, - 0x78,0x36,0x31,0x31,0x32,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x4b,0x41,0x57,0x50,0x4f,0x57,0x5f,0x52,0x56, - 0x4e,0x20,0x30,0x78,0x36,0x62,0x30,0x66,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x58,0x4c,0x41, - 0x20,0x30,0x78,0x37,0x32,0x31,0x32,0x31,0x31,0x66,0x66,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x41,0x4d,0x49,0x4c,0x59,0x5f,0x55,0x4e,0x4b,0x4e,0x4f, + 0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x58,0x4c,0x41,0x20,0x30,0x78,0x37,0x32,0x31,0x32,0x31,0x31,0x66,0x66,0x0a,0x23, + 0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x53,0x46,0x58,0x20,0x30,0x78,0x37,0x32,0x31,0x35,0x31,0x32,0x37,0x33,0x0a,0x23,0x64, + 0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x47,0x52,0x41,0x46,0x54,0x20,0x30,0x78,0x37,0x32,0x31,0x35,0x31,0x32,0x36,0x37,0x0a,0x23, + 0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x52,0x58,0x5f,0x59,0x41,0x44,0x41,0x20,0x30,0x78,0x37,0x32,0x31,0x35,0x31,0x32,0x37,0x39,0x0a,0x23, + 0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,0x41,0x20,0x30,0x78,0x36,0x31,0x31,0x33,0x30,0x30,0x30, + 0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x43,0x48,0x55,0x4b,0x57,0x41,0x5f,0x56,0x32,0x20,0x30,0x78,0x36, + 0x31,0x31,0x34,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x41,0x52,0x32,0x5f,0x57,0x52,0x4b,0x5a,0x20,0x30,0x78, + 0x36,0x31,0x31,0x32,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x41,0x4c,0x47,0x4f,0x5f,0x4b,0x41,0x57,0x50,0x4f,0x57,0x5f,0x52,0x56,0x4e, + 0x20,0x30,0x78,0x36,0x62,0x30,0x66,0x30,0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x41,0x4d,0x49,0x4c,0x59,0x5f,0x55,0x4e,0x4b,0x4e,0x4f, 0x57,0x4e,0x20,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x41,0x4d,0x49,0x4c,0x59,0x5f,0x43,0x4e,0x20,0x30,0x78,0x36,0x33,0x31,0x35,0x30,0x30,0x30, 0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x41,0x4d,0x49,0x4c,0x59,0x5f,0x43,0x4e,0x5f,0x4c,0x49,0x54,0x45,0x20,0x30,0x78,0x36,0x33,0x31,0x34,0x30, 0x30,0x30,0x30,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x41,0x4d,0x49,0x4c,0x59,0x5f,0x43,0x4e,0x5f,0x48,0x45,0x41,0x56,0x59,0x20,0x30,0x78,0x36,0x33, diff --git a/src/base/crypto/Algorithm.cpp b/src/base/crypto/Algorithm.cpp index c15da6d0ea..05bf2e9636 100644 --- a/src/base/crypto/Algorithm.cpp +++ b/src/base/crypto/Algorithm.cpp @@ -85,6 +85,7 @@ const char *Algorithm::kRX_V2 = "rx/2"; const char *Algorithm::kRX_WOW = "rx/wow"; const char *Algorithm::kRX_ARQ = "rx/arq"; const char *Algorithm::kRX_XEQ = "rx/xeq"; +const char *Algorithm::kRX_XLA = "panthera"; const char *Algorithm::kRX_GRAFT = "rx/graft"; const char *Algorithm::kRX_SFX = "rx/sfx"; const char *Algorithm::kRX_KEVA = "rx/keva"; @@ -110,10 +111,6 @@ const char* Algorithm::kFLEX = "flex"; const char* Algorithm::kFLEX_KCN = "flex"; #endif -#ifdef XMRIG_ALGO_RANDOMX -const char *Algorithm::kRX_XLA = "panthera"; -#endif - #define ALGO_NAME(ALGO) { Algorithm::ALGO, Algorithm::k##ALGO } #define ALGO_ALIAS(ALGO, NAME) { NAME, Algorithm::ALGO } @@ -291,6 +288,7 @@ static const std::map kAlgorithmAlias ALGO_ALIAS(RX_ARQ, "randomarq"), ALGO_ALIAS_AUTO(RX_XEQ), ALGO_ALIAS(RX_XEQ, "randomx/xeq"), ALGO_ALIAS(RX_XEQ, "randomxeq"), + ALGO_ALIAS_AUTO(RX_XLA), ALGO_ALIAS(RX_XLA, "Panthera"), ALGO_ALIAS_AUTO(RX_GRAFT), ALGO_ALIAS(RX_GRAFT, "randomx/graft"), ALGO_ALIAS(RX_GRAFT, "randomgraft"), ALGO_ALIAS_AUTO(RX_SFX), ALGO_ALIAS(RX_SFX, "randomx/sfx"), @@ -311,10 +309,6 @@ static const std::map kAlgorithmAlias ALGO_ALIAS_AUTO(KAWPOW_RVN), ALGO_ALIAS(KAWPOW_RVN, "kawpow/rvn"), # endif -# ifdef XMRIG_ALGO_RANDOMX - ALGO_ALIAS_AUTO(RX_XLA), ALGO_ALIAS(RX_XLA, "Panthera"), -# endif - # ifdef XMRIG_ALGO_GHOSTRIDER ALGO_ALIAS_AUTO(GHOSTRIDER_RTM), ALGO_ALIAS(GHOSTRIDER_RTM, "ghostrider/rtm"), ALGO_ALIAS(GHOSTRIDER_RTM, "gr"), diff --git a/src/base/crypto/Algorithm.h b/src/base/crypto/Algorithm.h index d8fc49462b..1c98e7cc33 100644 --- a/src/base/crypto/Algorithm.h +++ b/src/base/crypto/Algorithm.h @@ -79,6 +79,7 @@ class Algorithm RX_WOW = 0x72141177, // "rx/wow" RandomWOW (Wownero). RX_ARQ = 0x72121061, // "rx/arq" RandomARQ (Arqma). RX_XEQ = 0x72121000, + RX_XLA = 0x721211ff, // "panthera" Panthera (Scala). RX_GRAFT = 0x72151267, // "rx/graft" RandomGRAFT (Graft). RX_SFX = 0x72151273, // "rx/sfx" RandomSFX (Safex Cash). RX_KEVA = 0x7214116b, // "rx/keva" RandomKEVA (Keva). @@ -87,8 +88,6 @@ class Algorithm AR2_CHUKWA_V2 = 0x61140000, // "argon2/chukwav2" Argon2id (Chukwa v2). AR2_WRKZ = 0x61120000, // "argon2/wrkz" Argon2id (WRKZ) KAWPOW_RVN = 0x6b0f0000, // "kawpow/rvn" KawPow (RVN) - - RX_XLA = 0x721211ff, // "panthera" Panthera (Scala2). }; enum Family : uint32_t { diff --git a/src/base/crypto/Coin.cpp b/src/base/crypto/Coin.cpp index d7122d5c51..4f27ab8650 100644 --- a/src/base/crypto/Coin.cpp +++ b/src/base/crypto/Coin.cpp @@ -67,6 +67,7 @@ const char *Coin::kDisabled = "DISABLED_COIN"; const char *Coin::kField = "coin"; const char *Coin::kUnknown = "UNKNOWN_COIN"; + } /* namespace xmrig */ diff --git a/src/base/io/log/Tags.h b/src/base/io/log/Tags.h index 5156ed1aae..2ce4535bf9 100644 --- a/src/base/io/log/Tags.h +++ b/src/base/io/log/Tags.h @@ -58,13 +58,13 @@ class Tags static const char *opencl(); # endif -# ifdef XMRIG_FEATURE_MO_BENCHMARK - static const char *benchmark(); -# endif - # ifdef XMRIG_FEATURE_PROFILING static const char* profiler(); # endif + +# ifdef XMRIG_FEATURE_MO_BENCHMARK + static const char *benchmark(); +# endif }; diff --git a/src/base/kernel/config/BaseConfig.cpp b/src/base/kernel/config/BaseConfig.cpp index 5e92195056..4561c10828 100644 --- a/src/base/kernel/config/BaseConfig.cpp +++ b/src/base/kernel/config/BaseConfig.cpp @@ -93,9 +93,6 @@ bool xmrig::BaseConfig::read(const IJsonReader &reader, const char *fileName) m_autoSave = reader.getBool(kAutosave, m_autoSave); m_background = reader.getBool(kBackground, m_background); m_dryRun = reader.getBool(kDryRun, m_dryRun); -# ifdef XMRIG_FEATURE_MO_BENCHMARK - m_rebenchAlgo = reader.getBool(kRebenchAlgo, m_rebenchAlgo); -# endif m_syslog = reader.getBool(kSyslog, m_syslog); m_watch = reader.getBool(kWatch, m_watch); m_logFile = reader.getString(kLogFile); @@ -107,11 +104,13 @@ bool xmrig::BaseConfig::read(const IJsonReader &reader, const char *fileName) m_tls = reader.getValue(kTls); # endif - Log::setColors(reader.getBool(kColors, Log::isColors())); # ifdef XMRIG_FEATURE_MO_BENCHMARK + m_rebenchAlgo = reader.getBool(kRebenchAlgo, m_rebenchAlgo); m_benchAlgoTime = reader.getInt(kBenchAlgoTime, m_benchAlgoTime); m_algoMinTime = reader.getInt(kAlgoMinTime, m_algoMinTime); # endif + + Log::setColors(reader.getBool(kColors, Log::isColors())); setVerbose(reader.getValue(kVerbose)); const auto &api = reader.getObject(kApi); diff --git a/src/base/kernel/config/BaseTransform.cpp b/src/base/kernel/config/BaseTransform.cpp index 40b12a8316..a2a4898230 100644 --- a/src/base/kernel/config/BaseTransform.cpp +++ b/src/base/kernel/config/BaseTransform.cpp @@ -239,18 +239,20 @@ void xmrig::BaseTransform::transform(rapidjson::Document &doc, int key, const ch # endif case IConfig::RetriesKey: /* --retries */ -# ifdef XMRIG_FEATURE_MO_BENCHMARK - case IConfig::BenchAlgoTimeKey: /* --bench-algo-time */ - case IConfig::AlgoMinTimeKey: /* --algo-min-time */ -# endif case IConfig::RetryPauseKey: /* --retry-pause */ case IConfig::PrintTimeKey: /* --print-time */ +# ifdef XMRIG_FEATURE_HTTP case IConfig::HttpPort: /* --http-port */ - case IConfig::DonateLevelKey: /* --donate-level */ case IConfig::DaemonPollKey: /* --daemon-poll-interval */ case IConfig::DaemonJobTimeoutKey: /* --daemon-job-timeout */ - case IConfig::DnsTtlKey: /* --dns-ttl */ case IConfig::DaemonZMQPortKey: /* --daemon-zmq-port */ +# endif + case IConfig::DonateLevelKey: /* --donate-level */ + case IConfig::DnsTtlKey: /* --dns-ttl */ +# ifdef XMRIG_FEATURE_MO_BENCHMARK + case IConfig::BenchAlgoTimeKey: /* --bench-algo-time */ + case IConfig::AlgoMinTimeKey: /* --algo-min-time */ +# endif return transformUint64(doc, key, static_cast(strtol(arg, nullptr, 10))); case IConfig::BackgroundKey: /* --background */ @@ -265,14 +267,13 @@ void xmrig::BaseTransform::transform(rapidjson::Document &doc, int key, const ch case IConfig::HttpEnabledKey: /* --http-enabled */ case IConfig::DaemonKey: /* --daemon */ # endif -# ifdef XMRIG_FEATURE_MO_BENCHMARK - case IConfig::RebenchAlgoKey: /* --rebench-algo */ -# endif - case IConfig::PauseOnBatteryKey: /* --pause-on-battery */ case IConfig::SubmitToOriginKey: /* --submit-to-origin */ case IConfig::VerboseKey: /* --verbose */ case IConfig::DnsIPv4Key: /* --ipv4 */ case IConfig::DnsIPv6Key: /* --ipv6 */ +# ifdef XMRIG_FEATURE_MO_BENCHMARK + case IConfig::RebenchAlgoKey: /* --rebench-algo */ +# endif return transformBoolean(doc, key, true); case IConfig::ColorKey: /* --no-color */ @@ -333,17 +334,17 @@ void xmrig::BaseTransform::transformBoolean(rapidjson::Document &doc, int key, b case IConfig::NoTitleKey: /* --no-title */ return set(doc, BaseConfig::kTitle, enable); -# ifdef XMRIG_FEATURE_MO_BENCHMARK - case IConfig::RebenchAlgoKey: /* --rebench-algo */ - return set(doc, BaseConfig::kRebenchAlgo, enable); -# endif - case IConfig::DnsIPv4Key: /* --ipv4 */ return set(doc, DnsConfig::kField, DnsConfig::kIPv, 4); case IConfig::DnsIPv6Key: /* --ipv6 */ return set(doc, DnsConfig::kField, DnsConfig::kIPv, 6); +# ifdef XMRIG_FEATURE_MO_BENCHMARK + case IConfig::RebenchAlgoKey: /* --rebench-algo */ + return set(doc, BaseConfig::kRebenchAlgo, enable); +# endif + default: break; } diff --git a/src/config.json b/src/config.json index f7d52d3bc0..2e06aa565e 100644 --- a/src/config.json +++ b/src/config.json @@ -55,8 +55,8 @@ "nvml": true, "cn/0": false, "cn-lite/0": false, - "panthera": false, - "astrobwt": false + "astrobwt": false, + "panthera": false }, "donate-level": 1, "donate-over-proxy": 1, @@ -102,8 +102,8 @@ "user-agent": null, "verbose": 0, "watch": true, - "rebench-algo": false, - "bench-algo-time": 20, "pause-on-battery": false, - "pause-on-active": false + "pause-on-active": false, + "rebench-algo": false, + "bench-algo-time": 20 } diff --git a/src/core/MoBenchmark.cpp b/src/core/MoBenchmark.cpp index 5001a84b38..20b29bc8c2 100644 --- a/src/core/MoBenchmark.cpp +++ b/src/core/MoBenchmark.cpp @@ -256,8 +256,3 @@ uint64_t MoBenchmark::get_now() const { // get current time in ms } } // namespace xmrig - -const char *xmrig::bm_tag() -{ - return Tags::benchmark(); -} diff --git a/src/core/config/Config.cpp b/src/core/config/Config.cpp index 9028f3e04a..7dcb7de079 100644 --- a/src/core/config/Config.cpp +++ b/src/core/config/Config.cpp @@ -245,14 +245,14 @@ bool xmrig::Config::read(const IJsonReader &reader, const char *fileName) d_ptr->healthPrintTime = reader.getUint(kHealthPrintTime, d_ptr->healthPrintTime); # endif -# ifdef XMRIG_FEATURE_MO_BENCHMARK - m_benchmark.read(reader.getValue(kAlgoPerf)); -# endif - # ifdef XMRIG_FEATURE_DMI d_ptr->dmi = reader.getBool(kDMI, d_ptr->dmi); # endif +# ifdef XMRIG_FEATURE_MO_BENCHMARK + m_benchmark.read(reader.getValue(kAlgoPerf)); +# endif + return true; } @@ -313,6 +313,8 @@ void xmrig::Config::getJSON(rapidjson::Document &doc) const doc.AddMember(StringRef(kUserAgent), m_userAgent.toJSON(), allocator); doc.AddMember(StringRef(kVerbose), Log::verbose(), allocator); doc.AddMember(StringRef(kWatch), m_watch, allocator); + doc.AddMember(StringRef(kPauseOnBattery), isPauseOnBattery(), allocator); + doc.AddMember(StringRef(kPauseOnActive), (d_ptr->idleTime == 0U || d_ptr->idleTime == kIdleTime) ? Value(isPauseOnActive()) : Value(d_ptr->idleTime), allocator); # ifdef XMRIG_FEATURE_MO_BENCHMARK doc.AddMember(StringRef(kRebenchAlgo), isRebenchAlgo(), allocator); @@ -320,7 +322,4 @@ void xmrig::Config::getJSON(rapidjson::Document &doc) const doc.AddMember(StringRef(kAlgoMinTime), algoMinTime(), allocator); doc.AddMember(StringRef(kAlgoPerf), m_benchmark.toJSON(doc), allocator); # endif - - doc.AddMember(StringRef(kPauseOnBattery), isPauseOnBattery(), allocator); - doc.AddMember(StringRef(kPauseOnActive), (d_ptr->idleTime == 0U || d_ptr->idleTime == kIdleTime) ? Value(isPauseOnActive()) : Value(d_ptr->idleTime), allocator); } diff --git a/src/core/config/Config.h b/src/core/config/Config.h index c744a9310e..5759f6a97f 100644 --- a/src/core/config/Config.h +++ b/src/core/config/Config.h @@ -104,7 +104,7 @@ class Config : public BaseConfig void getJSON(rapidjson::Document &doc) const override; # ifdef XMRIG_FEATURE_MO_BENCHMARK - inline MoBenchmark &benchmark() { return m_benchmark; } + inline MoBenchmark &benchmark() { return m_benchmark; } # endif private: diff --git a/src/crypto/cn/CnAlgo.h b/src/crypto/cn/CnAlgo.h index ab4402795b..2f82b7ae42 100644 --- a/src/crypto/cn/CnAlgo.h +++ b/src/crypto/cn/CnAlgo.h @@ -40,6 +40,9 @@ class 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; } +# ifdef XMRIG_ALGO_CN_GPU + constexpr inline bool isGPU() const { return ALGO == Algorithm::CN_GPU; } +# endif 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; } constexpr inline uint32_t mask() const { return static_cast(((memory() - 1) / 16) * 16); } diff --git a/src/crypto/cn/CryptoNight_x86.h b/src/crypto/cn/CryptoNight_x86.h index 8e1486a94f..f0a281eafa 100644 --- a/src/crypto/cn/CryptoNight_x86.h +++ b/src/crypto/cn/CryptoNight_x86.h @@ -301,7 +301,7 @@ static NOINLINE void cn_explode_scratchpad(cryptonight_ctx *ctx) constexpr CnAlgo props; # ifdef XMRIG_ALGO_CN_GPU - constexpr bool IS_HEAVY = props.isHeavy() || ALGO == Algorithm::CN_GPU; + constexpr bool IS_HEAVY = props.isHeavy() || props.isGPU(); # else constexpr bool IS_HEAVY = props.isHeavy(); # endif @@ -321,6 +321,31 @@ static NOINLINE void cn_explode_scratchpad(cryptonight_ctx *ctx) const __m128i* input = reinterpret_cast(ctx->state); __m128i* output = reinterpret_cast<__m128i*>(ctx->memory); +# ifdef XMRIG_ALGO_CN_GPU + if (props.isGPU()) { + constexpr size_t hash_size = 200; // 25x8 bytes + alignas(16) uint64_t hash[25]; + + for (uint64_t i = 0; i < props.memory() / 512; i++) { + memcpy(hash, input, hash_size); + hash[0] ^= i; + + keccakf(hash, 24); + memcpy(output, hash, 160); + output += 160; + + keccakf(hash, 24); + memcpy(output, hash, 176); + output += 176; + + keccakf(hash, 24); + memcpy(output, hash, 176); + output += 176; + } + return; + } +# endif + aes_genkey(input, &k0, &k1, &k2, &k3, &k4, &k5, &k6, &k7, &k8, &k9); if (props.half_mem() && !ctx->first_half) { @@ -421,7 +446,7 @@ static NOINLINE void cn_implode_scratchpad(cryptonight_ctx *ctx) constexpr CnAlgo props; # ifdef XMRIG_ALGO_CN_GPU - constexpr bool IS_HEAVY = props.isHeavy() || ALGO == Algorithm::CN_GPU; + constexpr bool IS_HEAVY = props.isHeavy() || props.isGPU(); # else constexpr bool IS_HEAVY = props.isHeavy(); # endif @@ -730,19 +755,207 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si VARIANT2_SET_ROUNDING_MODE(); VARIANT4_RANDOM_MATH_INIT(0); - uint64_t al0 = h0[0] ^ h0[4]; - uint64_t ah0 = h0[1] ^ h0[5]; - uint64_t idx0 = al0; - __m128i bx0 = _mm_set_epi64x(static_cast(h0[3] ^ h0[7]), static_cast(h0[2] ^ h0[6])); - __m128i bx1 = _mm_set_epi64x(static_cast(h0[9] ^ h0[11]), static_cast(h0[8] ^ h0[10])); - + uint64_t al0; + uint64_t ah0; + uint64_t idx0; + __m128i bx0; + __m128i bx1; __m128 conc_var; +# ifdef XMRIG_ALGO_CN_GPU + constexpr bool isAVX2 = Cpu::info()->hasAVX2(); + const uint8_t* spad = ctx[0]->state; + uint8_t* lpad = ctx[0]->memory; + +# ifndef _mm256_bslli_epi128 +# define _mm256_bslli_epi128(a, count) _mm256_slli_si256((a), (count)) +# endif +# ifndef _mm256_bsrli_epi128 +# define _mm256_bsrli_epi128(a, count) _mm256_srli_si256((a), (count)) +# endif + + uint32_t s; + __m128i* idx0s; + __m128i* idx1s; + __m128i* idx2s; + __m128i* idx3s; + __m128 sum0s; + __m128 n0, n1, n2, n3; + __m128i v0, v1, v2, v3; + __m128 sumas, sumbs, sum1s, sum2s, sum3s; + + __m256i* idx0a; + __m256i* idx2a; + __m256 sum0a; + __m256i v01, v23; + __m256 sumaa, sumba, sum1a; + __m256 rc = sum0a; + __m256 n01, n23; + __m256i out, out2; + __m256 n10, n22, n33; + __m256 n11, n02, n30; + __m128 sum; + + if (!props.isGPU()) { +# endif + al0 = h0[0] ^ h0[4]; + ah0 = h0[1] ^ h0[5]; + idx0 = al0; + bx0 = _mm_set_epi64x(static_cast(h0[3] ^ h0[7]), static_cast(h0[2] ^ h0[6])); + bx1 = _mm_set_epi64x(static_cast(h0[9] ^ h0[11]), static_cast(h0[8] ^ h0[10])); + if (ALGO == Algorithm::CN_CCX) { conc_var = _mm_setzero_ps(); RESTORE_ROUNDING_MODE(); } +# ifdef XMRIG_ALGO_CN_GPU + } else { + RESTORE_ROUNDING_MODE(); + s = reinterpret_cast(spad)[0] >> 8; + if (isAVX2) { + idx0a = scratchpad_ptr(lpad, s, 0); + idx2a = scratchpad_ptr(lpad, s, 2); + sum0a = _mm256_setzero_ps(); + } else { + idx0s = scratchpad_ptr(lpad, s, 0); + idx1s = scratchpad_ptr(lpad, s, 1); + idx2s = scratchpad_ptr(lpad, s, 2); + idx3s = scratchpad_ptr(lpad, s, 3); + sum0s = _mm_setzero_ps(); + } + } + +# endif for (size_t i = 0; i < props.iterations(); i++) { +# ifdef XMRIG_ALGO_CN_GPU + if (props.isGPU()) { + if (isAVX2) { + // cn_gpu_inner_avx(ctx[0]->state, ctx[0]->memory); + rc = sum0a; + + prep_dv_avx(idx0a, v01, n01); + prep_dv_avx(idx2a, v23, n23); + + n10 = _mm256_permute2f128_ps(n01, n01, 0x01); + n22 = _mm256_permute2f128_ps(n23, n23, 0x00); + n33 = _mm256_permute2f128_ps(n23, n23, 0x11); + + out = _mm256_setzero_si256(); + double_compute_wrap<0>(n01, n10, n22, n33, 1.3437500f, 1.4296875f, rc, suma, out); + double_compute_wrap<1>(n01, n22, n33, n10, 1.2812500f, 1.3984375f, rc, suma, out); + double_compute_wrap<2>(n01, n33, n10, n22, 1.3593750f, 1.3828125f, rc, sumb, out); + double_compute_wrap<3>(n01, n33, n22, n10, 1.3671875f, 1.3046875f, rc, sumb, out); + _mm256_store_si256(idx0a, _mm256_xor_si256(v01, out)); + sum0a = _mm256_add_ps(suma, sumb); + out2 = out; + + n11 = _mm256_permute2f128_ps(n01, n01, 0x11); + n02 = _mm256_permute2f128_ps(n01, n23, 0x20); + n30 = _mm256_permute2f128_ps(n01, n23, 0x03); + + out = _mm256_setzero_si256(); + double_compute_wrap<0>(n23, n11, n02, n30, 1.4140625f, 1.3203125f, rc, suma, out); + double_compute_wrap<1>(n23, n02, n30, n11, 1.2734375f, 1.3515625f, rc, suma, out); + double_compute_wrap<2>(n23, n30, n11, n02, 1.2578125f, 1.3359375f, rc, sumb, out); + double_compute_wrap<3>(n23, n30, n02, n11, 1.2890625f, 1.4609375f, rc, sumb, out); + _mm256_store_si256(idx2a, _mm256_xor_si256(v23, out)); + sum1 = _mm256_add_ps(suma, sumb); + + out2 = _mm256_xor_si256(out2, out); + out2 = _mm256_xor_si256(_mm256_permute2x128_si256(out2,out2,0x41), out2); + suma = _mm256_permute2f128_ps(sum0a, sum1, 0x30); + sumb = _mm256_permute2f128_ps(sum0a, sum1, 0x21); + sum0a = _mm256_add_ps(suma, sumb); + sum0a = _mm256_add_ps(sum0a, _mm256_permute2f128_ps(sum0a, sum0a, 0x41)); + + // Clear the high 128 bits + sum = _mm256_castps256_ps128(sum0a); + + sum = _mm_and_ps(_mm_castsi128_ps(_mm_set1_epi32(0x7fffffff)), sum); // take abs(va) by masking the float sign bit + // vs range 0 - 64 + __m128i v0 = _mm_cvttps_epi32(_mm_mul_ps(sum, _mm_set1_ps(16777216.0f))); + v0 = _mm_xor_si128(v0, _mm256_castsi256_si128(out2)); + __m128i v1 = _mm_shuffle_epi32(v0, _MM_SHUFFLE(0, 1, 2, 3)); + v0 = _mm_xor_si128(v0, v1); + v1 = _mm_shuffle_epi32(v0, _MM_SHUFFLE(0, 1, 0, 1)); + v0 = _mm_xor_si128(v0, v1); + + // vs is now between 0 and 1 + sum = _mm_div_ps(sum, _mm_set1_ps(64.0f)); + sum0a = _mm256_insertf128_ps(_mm256_castps128_ps256(sum), sum, 1); + uint32_t n = _mm_cvtsi128_si32(v0); + idx0a = scratchpad_ptr(lpad, n, 0); + idx2a = scratchpad_ptr(lpad, n, 2); + } else { + //cn_gpu_inner_ssse3(ctx[0]->state, ctx[0]->memory); + prep_dv(idx0s, v0, n0); + prep_dv(idx1s, v1, n1); + prep_dv(idx2s, v2, n2); + prep_dv(idx3s, v3, n3); + __m128 rc = sum0s; + + __m128i out, out2; + out = _mm_setzero_si128(); + single_compute_wrap<0>(n0, n1, n2, n3, 1.3437500f, rc, suma, out); + single_compute_wrap<1>(n0, n2, n3, n1, 1.2812500f, rc, suma, out); + single_compute_wrap<2>(n0, n3, n1, n2, 1.3593750f, rc, sumb, out); + single_compute_wrap<3>(n0, n3, n2, n1, 1.3671875f, rc, sumb, out); + sum0s = _mm_add_ps(suma, sumb); + _mm_store_si128(idx0s, _mm_xor_si128(v0, out)); + out2 = out; + + out = _mm_setzero_si128(); + single_compute_wrap<0>(n1, n0, n2, n3, 1.4296875f, rc, suma, out); + single_compute_wrap<1>(n1, n2, n3, n0, 1.3984375f, rc, suma, out); + single_compute_wrap<2>(n1, n3, n0, n2, 1.3828125f, rc, sumb, out); + single_compute_wrap<3>(n1, n3, n2, n0, 1.3046875f, rc, sumb, out); + sum1 = _mm_add_ps(suma, sumb); + _mm_store_si128(idx1s, _mm_xor_si128(v1, out)); + out2 = _mm_xor_si128(out2, out); + + out = _mm_setzero_si128(); + single_compute_wrap<0>(n2, n1, n0, n3, 1.4140625f, rc, suma, out); + single_compute_wrap<1>(n2, n0, n3, n1, 1.2734375f, rc, suma, out); + single_compute_wrap<2>(n2, n3, n1, n0, 1.2578125f, rc, sumb, out); + single_compute_wrap<3>(n2, n3, n0, n1, 1.2890625f, rc, sumb, out); + sum2 = _mm_add_ps(suma, sumb); + _mm_store_si128(idx2s, _mm_xor_si128(v2, out)); + out2 = _mm_xor_si128(out2, out); + + out = _mm_setzero_si128(); + single_compute_wrap<0>(n3, n1, n2, n0, 1.3203125f, rc, suma, out); + single_compute_wrap<1>(n3, n2, n0, n1, 1.3515625f, rc, suma, out); + single_compute_wrap<2>(n3, n0, n1, n2, 1.3359375f, rc, sumb, out); + single_compute_wrap<3>(n3, n0, n2, n1, 1.4609375f, rc, sumb, out); + sum3 = _mm_add_ps(suma, sumb); + _mm_store_si128(idx3s, _mm_xor_si128(v3, out)); + out2 = _mm_xor_si128(out2, out); + sum0s = _mm_add_ps(sum0s, sum1); + sum2 = _mm_add_ps(sum2, sum3); + sum0s = _mm_add_ps(sum0s, sum2); + + sum0s = _mm_and_ps(_mm_castsi128_ps(_mm_set1_epi32(0x7fffffff)), sum0s); // take abs(va) by masking the float sign bit + // vs range 0 - 64 + n0 = _mm_mul_ps(sum0s, _mm_set1_ps(16777216.0f)); + v0 = _mm_cvttps_epi32(n0); + v0 = _mm_xor_si128(v0, out2); + v1 = _mm_shuffle_epi32(v0, _MM_SHUFFLE(0, 1, 2, 3)); + v0 = _mm_xor_si128(v0, v1); + v1 = _mm_shuffle_epi32(v0, _MM_SHUFFLE(0, 1, 0, 1)); + v0 = _mm_xor_si128(v0, v1); + + // vs is now between 0 and 1 + sum0s = _mm_div_ps(sum0s, _mm_set1_ps(64.0f)); + uint32_t n = _mm_cvtsi128_si32(v0); + idx0s = scratchpad_ptr(lpad, n, 0); + idx1s = scratchpad_ptr(lpad, n, 1); + idx2s = scratchpad_ptr(lpad, n, 2); + idx3s = scratchpad_ptr(lpad, n, 3); + } + continue; + } +# endif + __m128i cx; if (IS_CN_HEAVY_TUBE || !SOFT_AES) { cx = _mm_load_si128(reinterpret_cast(&l0[interleaved_index(idx0 & MASK)])); @@ -859,6 +1072,9 @@ inline void cryptonight_single_hash(const uint8_t *__restrict__ input, size_t si cn_implode_scratchpad(ctx[0]); keccakf(h0, 24); +# ifdef XMRIG_ALGO_CN_GPU + if (props.isGPU()) { memcpy(output, ctx[0]->state, 32); return; } +# endif if (height == 101) // Flex algo ugly hack extra_hashes_flex[ctx[0]->state[0] & 2](ctx[0]->state, 200, output); else @@ -881,41 +1097,15 @@ void cn_gpu_inner_ssse3(const uint8_t *spad, uint8_t *lpad); namespace xmrig { -template -static NOINLINE void cn_explode_scratchpad_gpu(cryptonight_ctx *ctx) -{ - const uint8_t* input = reinterpret_cast(ctx->state); - uint8_t* output = reinterpret_cast(ctx->memory); - - constexpr size_t hash_size = 200; // 25x8 bytes - alignas(16) uint64_t hash[25]; - - for (uint64_t i = 0; i < MEM / 512; i++) { - memcpy(hash, input, hash_size); - hash[0] ^= i; - - xmrig::keccakf(hash, 24); - memcpy(output, hash, 160); - output += 160; - - xmrig::keccakf(hash, 24); - memcpy(output, hash, 176); - output += 176; - - xmrig::keccakf(hash, 24); - memcpy(output, hash, 176); - output += 176; - } -} - - template inline void cryptonight_single_hash_gpu(const uint8_t *__restrict__ input, size_t size, uint8_t *__restrict__ output, cryptonight_ctx **__restrict__ ctx, uint64_t height) { constexpr CnAlgo props; + constexpr size_t MASK = props.mask(); + constexpr Algorithm::Id BASE = props.base(); keccak(input, size, ctx[0]->state); - cn_explode_scratchpad_gpu(ctx[0]); + cn_explode_scratchpad(ctx[0]); # ifdef _MSC_VER _control87(RC_NEAR, MCW_RC); @@ -924,12 +1114,12 @@ inline void cryptonight_single_hash_gpu(const uint8_t *__restrict__ input, size_ # endif if (xmrig::Cpu::info()->hasAVX2()) { - cn_gpu_inner_avx(ctx[0]->state, ctx[0]->memory); + cn_gpu_inner_avx(ctx[0]->state, ctx[0]->memory); } else { - cn_gpu_inner_ssse3(ctx[0]->state, ctx[0]->memory); + cn_gpu_inner_ssse3(ctx[0]->state, ctx[0]->memory); } - cn_implode_scratchpad(ctx[0]); + cn_implode_scratchpad(ctx[0]); keccakf(reinterpret_cast(ctx[0]->state), 24); memcpy(output, ctx[0]->state, 32); } diff --git a/src/crypto/cn/gpu/cn_gpu_avx.cpp b/src/crypto/cn/gpu/cn_gpu_avx.cpp index c918196f8a..fae59f97ab 100644 --- a/src/crypto/cn/gpu/cn_gpu_avx.cpp +++ b/src/crypto/cn/gpu/cn_gpu_avx.cpp @@ -22,29 +22,32 @@ * along with this program. If not, see . */ - -#include "crypto/cn/CnAlgo.h" - - #ifdef __GNUC__ # include #else # include # define __restrict__ __restrict #endif + + +#include "crypto/cn/CnAlgo.h" + + #ifndef _mm256_bslli_epi128 - #define _mm256_bslli_epi128(a, count) _mm256_slli_si256((a), (count)) +# define _mm256_bslli_epi128(a, count) _mm256_slli_si256((a), (count)) #endif #ifndef _mm256_bsrli_epi128 - #define _mm256_bsrli_epi128(a, count) _mm256_srli_si256((a), (count)) +# define _mm256_bsrli_epi128(a, count) _mm256_srli_si256((a), (count)) #endif + inline void prep_dv_avx(__m256i* idx, __m256i& v, __m256& n01) { v = _mm256_load_si256(idx); n01 = _mm256_cvtepi32_ps(v); } + inline __m256 fma_break(const __m256& x) { // Break the dependency chain by setting the exp to ?????01 @@ -52,6 +55,7 @@ inline __m256 fma_break(const __m256& x) return _mm256_or_ps(_mm256_castsi256_ps(_mm256_set1_epi32(0x00800000)), xx); } + // 14 inline void sub_round(const __m256& n0, const __m256& n1, const __m256& n2, const __m256& n3, const __m256& rnd_c, __m256& n, __m256& d, __m256& c) { @@ -138,19 +142,18 @@ template void cn_gpu_inner_avx(const uint8_t* spad, uint8_t* lpad) { uint32_t s = reinterpret_cast(spad)[0] >> 8; - __m256i* idx0 = scratchpad_ptr(lpad, s, 0); - __m256i* idx2 = scratchpad_ptr(lpad, s, 2); + __m256i* gdx0 = scratchpad_ptr(lpad, s, 0); + __m256i* gdx2 = scratchpad_ptr(lpad, s, 2); __m256 sum0 = _mm256_setzero_ps(); - for(size_t i = 0; i < ITER; i++) - { + for (size_t i = 0; i < ITER; i++) { __m256i v01, v23; __m256 suma, sumb, sum1; __m256 rc = sum0; __m256 n01, n23; - prep_dv_avx(idx0, v01, n01); - prep_dv_avx(idx2, v23, n23); + prep_dv_avx(gdx0, v01, n01); + prep_dv_avx(gdx2, v23, n23); __m256i out, out2; __m256 n10, n22, n33; @@ -163,7 +166,7 @@ void cn_gpu_inner_avx(const uint8_t* spad, uint8_t* lpad) double_compute_wrap<1>(n01, n22, n33, n10, 1.2812500f, 1.3984375f, rc, suma, out); double_compute_wrap<2>(n01, n33, n10, n22, 1.3593750f, 1.3828125f, rc, sumb, out); double_compute_wrap<3>(n01, n33, n22, n10, 1.3671875f, 1.3046875f, rc, sumb, out); - _mm256_store_si256(idx0, _mm256_xor_si256(v01, out)); + _mm256_store_si256(gdx0, _mm256_xor_si256(v01, out)); sum0 = _mm256_add_ps(suma, sumb); out2 = out; @@ -177,7 +180,7 @@ void cn_gpu_inner_avx(const uint8_t* spad, uint8_t* lpad) double_compute_wrap<1>(n23, n02, n30, n11, 1.2734375f, 1.3515625f, rc, suma, out); double_compute_wrap<2>(n23, n30, n11, n02, 1.2578125f, 1.3359375f, rc, sumb, out); double_compute_wrap<3>(n23, n30, n02, n11, 1.2890625f, 1.4609375f, rc, sumb, out); - _mm256_store_si256(idx2, _mm256_xor_si256(v23, out)); + _mm256_store_si256(gdx2, _mm256_xor_si256(v23, out)); sum1 = _mm256_add_ps(suma, sumb); out2 = _mm256_xor_si256(out2, out); @@ -203,8 +206,8 @@ void cn_gpu_inner_avx(const uint8_t* spad, uint8_t* lpad) sum = _mm_div_ps(sum, _mm_set1_ps(64.0f)); sum0 = _mm256_insertf128_ps(_mm256_castps128_ps256(sum), sum, 1); uint32_t n = _mm_cvtsi128_si32(v0); - idx0 = scratchpad_ptr(lpad, n, 0); - idx2 = scratchpad_ptr(lpad, n, 2); + gdx0 = scratchpad_ptr(lpad, n, 0); + gdx2 = scratchpad_ptr(lpad, n, 2); } } diff --git a/src/crypto/cn/gpu/cn_gpu_ssse3.cpp b/src/crypto/cn/gpu/cn_gpu_ssse3.cpp index 658fe86a7a..83e4532d62 100644 --- a/src/crypto/cn/gpu/cn_gpu_ssse3.cpp +++ b/src/crypto/cn/gpu/cn_gpu_ssse3.cpp @@ -22,10 +22,6 @@ * along with this program. If not, see . */ - -#include "crypto/cn/CnAlgo.h" - - #ifdef __GNUC__ # include #else @@ -33,6 +29,10 @@ # define __restrict__ __restrict #endif + +#include "crypto/cn/CnAlgo.h" + + inline void prep_dv(__m128i* idx, __m128i& v, __m128& n) { v = _mm_load_si128(idx); @@ -131,22 +131,21 @@ template void cn_gpu_inner_ssse3(const uint8_t* spad, uint8_t* lpad) { uint32_t s = reinterpret_cast(spad)[0] >> 8; - __m128i* idx0 = scratchpad_ptr(lpad, s, 0); - __m128i* idx1 = scratchpad_ptr(lpad, s, 1); - __m128i* idx2 = scratchpad_ptr(lpad, s, 2); - __m128i* idx3 = scratchpad_ptr(lpad, s, 3); + __m128i* gdx0 = scratchpad_ptr(lpad, s, 0); + __m128i* gdx1 = scratchpad_ptr(lpad, s, 1); + __m128i* gdx2 = scratchpad_ptr(lpad, s, 2); + __m128i* gdx3 = scratchpad_ptr(lpad, s, 3); __m128 sum0 = _mm_setzero_ps(); - for(size_t i = 0; i < ITER; i++) - { + for (size_t i = 0; i < ITER; i++) { __m128 n0, n1, n2, n3; __m128i v0, v1, v2, v3; __m128 suma, sumb, sum1, sum2, sum3; - prep_dv(idx0, v0, n0); - prep_dv(idx1, v1, n1); - prep_dv(idx2, v2, n2); - prep_dv(idx3, v3, n3); + prep_dv(gdx0, v0, n0); + prep_dv(gdx1, v1, n1); + prep_dv(gdx2, v2, n2); + prep_dv(gdx3, v3, n3); __m128 rc = sum0; __m128i out, out2; @@ -156,7 +155,7 @@ void cn_gpu_inner_ssse3(const uint8_t* spad, uint8_t* lpad) single_compute_wrap<2>(n0, n3, n1, n2, 1.3593750f, rc, sumb, out); single_compute_wrap<3>(n0, n3, n2, n1, 1.3671875f, rc, sumb, out); sum0 = _mm_add_ps(suma, sumb); - _mm_store_si128(idx0, _mm_xor_si128(v0, out)); + _mm_store_si128(gdx0, _mm_xor_si128(v0, out)); out2 = out; out = _mm_setzero_si128(); @@ -165,7 +164,7 @@ void cn_gpu_inner_ssse3(const uint8_t* spad, uint8_t* lpad) single_compute_wrap<2>(n1, n3, n0, n2, 1.3828125f, rc, sumb, out); single_compute_wrap<3>(n1, n3, n2, n0, 1.3046875f, rc, sumb, out); sum1 = _mm_add_ps(suma, sumb); - _mm_store_si128(idx1, _mm_xor_si128(v1, out)); + _mm_store_si128(gdx1, _mm_xor_si128(v1, out)); out2 = _mm_xor_si128(out2, out); out = _mm_setzero_si128(); @@ -174,7 +173,7 @@ void cn_gpu_inner_ssse3(const uint8_t* spad, uint8_t* lpad) single_compute_wrap<2>(n2, n3, n1, n0, 1.2578125f, rc, sumb, out); single_compute_wrap<3>(n2, n3, n0, n1, 1.2890625f, rc, sumb, out); sum2 = _mm_add_ps(suma, sumb); - _mm_store_si128(idx2, _mm_xor_si128(v2, out)); + _mm_store_si128(gdx2, _mm_xor_si128(v2, out)); out2 = _mm_xor_si128(out2, out); out = _mm_setzero_si128(); @@ -183,7 +182,7 @@ void cn_gpu_inner_ssse3(const uint8_t* spad, uint8_t* lpad) single_compute_wrap<2>(n3, n0, n1, n2, 1.3359375f, rc, sumb, out); single_compute_wrap<3>(n3, n0, n2, n1, 1.4609375f, rc, sumb, out); sum3 = _mm_add_ps(suma, sumb); - _mm_store_si128(idx3, _mm_xor_si128(v3, out)); + _mm_store_si128(gdx3, _mm_xor_si128(v3, out)); out2 = _mm_xor_si128(out2, out); sum0 = _mm_add_ps(sum0, sum1); sum2 = _mm_add_ps(sum2, sum3); @@ -202,10 +201,10 @@ void cn_gpu_inner_ssse3(const uint8_t* spad, uint8_t* lpad) // vs is now between 0 and 1 sum0 = _mm_div_ps(sum0, _mm_set1_ps(64.0f)); uint32_t n = _mm_cvtsi128_si32(v0); - idx0 = scratchpad_ptr(lpad, n, 0); - idx1 = scratchpad_ptr(lpad, n, 1); - idx2 = scratchpad_ptr(lpad, n, 2); - idx3 = scratchpad_ptr(lpad, n, 3); + gdx0 = scratchpad_ptr(lpad, n, 0); + gdx1 = scratchpad_ptr(lpad, n, 1); + gdx2 = scratchpad_ptr(lpad, n, 2); + gdx3 = scratchpad_ptr(lpad, n, 3); } } diff --git a/src/crypto/randomx/common.hpp b/src/crypto/randomx/common.hpp index 86f073cf70..0f4ca14801 100644 --- a/src/crypto/randomx/common.hpp +++ b/src/crypto/randomx/common.hpp @@ -74,7 +74,7 @@ namespace randomx { constexpr int SuperscalarMaxSize = 3 * RANDOMX_SUPERSCALAR_MAX_LATENCY + 2; constexpr size_t CacheLineSize = RANDOMX_DATASET_ITEM_SIZE; #define ScratchpadSize RandomX_CurrentConfig.ScratchpadL3_Size - #define CacheLineAlignMask RandomX_CurrentConfig.CacheLineAlignMask_Calculated + #define CacheLineAlignMask RandomX_CurrentConfig.CacheLineAlignMask_Calculated #define DatasetExtraItems RandomX_ConfigurationBase::DatasetExtraItems_Calculated constexpr int StoreL3Condition = 14; diff --git a/src/crypto/randomx/panthera/sysendian.h b/src/crypto/randomx/panthera/sysendian.h index 52c1fe73b2..c51730d1cd 100644 --- a/src/crypto/randomx/panthera/sysendian.h +++ b/src/crypto/randomx/panthera/sysendian.h @@ -35,6 +35,8 @@ #define be64enc libcperciva_be64enc #define le32dec libcperciva_le32dec #define le32enc libcperciva_le32enc +#define le64dec libcperciva_le64dec +#define le64enc libcperciva_le64enc static inline uint32_t be32dec(const void * pp) @@ -91,4 +93,30 @@ le32enc(void * pp, uint32_t x) p[3] = (x >> 24) & 0xff; } +static inline uint64_t +le64dec(const void * pp) +{ + const uint8_t * p = (uint8_t const *)pp; + + return ((uint64_t)(p[0]) + ((uint64_t)(p[1]) << 8) + + ((uint64_t)(p[2]) << 16) + ((uint64_t)(p[3]) << 24) + + ((uint64_t)(p[4]) << 32) + ((uint64_t)(p[5]) << 40) + + ((uint64_t)(p[6]) << 48) + ((uint64_t)(p[7]) << 56)); +} + +static inline void +le64enc(void * pp, uint64_t x) +{ + uint8_t * p = (uint8_t *)pp; + + p[0] = x & 0xff; + p[1] = (x >> 8) & 0xff; + p[2] = (x >> 16) & 0xff; + p[3] = (x >> 24) & 0xff; + p[4] = (x >> 32) & 0xff; + p[5] = (x >> 40) & 0xff; + p[6] = (x >> 48) & 0xff; + p[7] = (x >> 56) & 0xff; +} + #endif /* !_SYSENDIAN_H_ */ diff --git a/src/crypto/randomx/randomx.cpp b/src/crypto/randomx/randomx.cpp index d9615f1fbf..68f18cff66 100644 --- a/src/crypto/randomx/randomx.cpp +++ b/src/crypto/randomx/randomx.cpp @@ -285,7 +285,7 @@ void RandomX_ConfigurationBase::Apply() ScratchpadL3Mask_Calculated = (((ScratchpadL3_Size / sizeof(uint64_t)) - 1) * 8); ScratchpadL3Mask64_Calculated = ((ScratchpadL3_Size / sizeof(uint64_t)) / 8 - 1) * 64; - CacheLineAlignMask_Calculated = (DatasetBaseSize - 1) & ~(RANDOMX_DATASET_ITEM_SIZE - 1); + CacheLineAlignMask_Calculated = (DatasetBaseSize - 1) & ~(RANDOMX_DATASET_ITEM_SIZE - 1); #if defined(XMRIG_FEATURE_ASM) && (defined(_M_X64) || defined(__x86_64__)) *(uint32_t*)(codeSshPrefetchTweaked + 3) = ArgonMemory * 16 - 1; @@ -668,8 +668,8 @@ extern "C" { assert(inputSize == 0 || input != nullptr); assert(output != nullptr); alignas(16) uint64_t tempHash[8]; - switch (algo) { - case xmrig::Algorithm::RX_XLA: rx_yespower_k12(tempHash, sizeof(tempHash), input, inputSize); break; + switch (algo) { + case xmrig::Algorithm::RX_XLA: rx_yespower_k12(tempHash, sizeof(tempHash), input, inputSize); break; default: rx_blake2b_wrapper::run(tempHash, sizeof(tempHash), input, inputSize); } machine->initScratchpad(&tempHash); @@ -683,8 +683,8 @@ extern "C" { } void randomx_calculate_hash_first(randomx_vm* machine, uint64_t (&tempHash)[8], const void* input, size_t inputSize, const xmrig::Algorithm algo) { - switch (algo) { - case xmrig::Algorithm::RX_XLA: rx_yespower_k12(tempHash, sizeof(tempHash), input, inputSize); break; + switch (algo) { + case xmrig::Algorithm::RX_XLA: rx_yespower_k12(tempHash, sizeof(tempHash), input, inputSize); break; default: rx_blake2b_wrapper::run(tempHash, sizeof(tempHash), input, inputSize); } machine->initScratchpad(tempHash); @@ -701,8 +701,8 @@ extern "C" { machine->run(&tempHash); // Finish current hash and fill the scratchpad for the next hash at the same time - switch (algo) { - case xmrig::Algorithm::RX_XLA: rx_yespower_k12(tempHash, sizeof(tempHash), nextInput, nextInputSize); break; + switch (algo) { + case xmrig::Algorithm::RX_XLA: rx_yespower_k12(tempHash, sizeof(tempHash), nextInput, nextInputSize); break; default: rx_blake2b_wrapper::run(tempHash, sizeof(tempHash), nextInput, nextInputSize); } machine->hashAndFill(output, tempHash); diff --git a/src/crypto/randomx/randomx.h b/src/crypto/randomx/randomx.h index c35f84b0ca..997dd7f928 100644 --- a/src/crypto/randomx/randomx.h +++ b/src/crypto/randomx/randomx.h @@ -76,9 +76,9 @@ struct RandomX_ConfigurationBase ConditionMask_Calculated = ((1 << JumpBits) - 1) << JumpOffset, }; - uint32_t ArgonMemory; - uint32_t CacheAccesses; - uint32_t DatasetBaseSize; + uint32_t ArgonMemory; + uint32_t CacheAccesses; + uint32_t DatasetBaseSize; uint32_t ArgonIterations; uint32_t ArgonLanes; @@ -139,11 +139,10 @@ struct RandomX_ConfigurationBase uint8_t codePrefetchScratchpadTweaked[28]; uint32_t codePrefetchScratchpadTweakedSize; - uint32_t CacheLineAlignMask_Calculated; - uint32_t AddressMask_Calculated[4]; uint32_t ScratchpadL3Mask_Calculated; uint32_t ScratchpadL3Mask64_Calculated; + uint32_t CacheLineAlignMask_Calculated; # if (XMRIG_ARM == 8) || defined(XMRIG_RISCV) uint32_t Log2_ScratchpadL1; diff --git a/src/version.h b/src/version.h index 92a8d6f789..f185a874e5 100644 --- a/src/version.h +++ b/src/version.h @@ -11,7 +11,7 @@ #define APP_ID "xmrig" #define APP_NAME "XMRig" #define APP_DESC "XMRig miner" -#define APP_VERSION "6.26.0-mo2" +#define APP_VERSION "6.26.0-mo2a" #define APP_DOMAIN "xmrig.com" #define APP_SITE "www.xmrig.com" #define APP_COPYRIGHT "Copyright (C) 2016-2026 xmrig.com"