Skip to content

Commit

Permalink
Merge xmrig-cuda v6.12.0 into master
Browse files Browse the repository at this point in the history
  • Loading branch information
@MoneroOcean
MoneroOcean committed Apr 20, 2021
2 parents e9b7953 + 37d72c8 commit 58a437c
Show file tree
Hide file tree
Showing 7 changed files with 90 additions and 16 deletions.
3 changes: 3 additions & 0 deletions CHANGELOG.md
View file
Original file line number Diff line number Diff line change
@@ -1,3 +1,6 @@
# v6.12.0
- [#95](https://github.com/xmrig/xmrig-cuda/pull/95) Added support for Uplexa (`cn/upx2` algorithm).

# v6.5.0
- [#74](https://github.com/xmrig/xmrig-cuda/pull/74) Fixed CUDA 8.0 support, RandomX, AstroBWT, and KawPow disabled for this CUDA version.
- [#76](https://github.com/xmrig/xmrig-cuda/pull/76) Fixed high CPU usage on Cryptonight and AstroBWT.
Expand Down
5 changes: 5 additions & 0 deletions CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -11,6 +11,7 @@ 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_ARGON2 "Enable Argon2 algorithms family" OFF) #unsupported

if (CUDA_VERSION VERSION_LESS 9.0)
Expand Down Expand Up @@ -40,6 +41,10 @@ if (WITH_CN_GPU)
add_definitions(/DXMRIG_ALGO_CN_GPU)
endif()

if (WITH_CN_FEMTO)
add_definitions(/DXMRIG_ALGO_CN_FEMTO)
endif()

if (WITH_RANDOMX)
add_definitions(/DXMRIG_ALGO_RANDOMX)
endif()
Expand Down
51 changes: 48 additions & 3 deletions src/crypto/cn/CnAlgo.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -66,6 +66,9 @@ class CnAlgo
case Algorithm::CN_PICO:
return CN_MEMORY / 8;

case Algorithm::CN_FEMTO:
return CN_MEMORY / 16;

default:
break;
}
Expand All @@ -85,11 +88,15 @@ class CnAlgo

case Algorithm::CN_FAST:
case Algorithm::CN_HALF:
# ifdef XMRIG_ALGO_CN_LITE
case Algorithm::CN_LITE_0:
case Algorithm::CN_LITE_1:
# endif
# ifdef XMRIG_ALGO_CN_HEAVY
case Algorithm::CN_HEAVY_0:
case Algorithm::CN_HEAVY_TUBE:
case Algorithm::CN_HEAVY_XHV:
# endif
case Algorithm::CN_CCX:
return CN_ITER / 2;

Expand All @@ -101,12 +108,21 @@ class CnAlgo
case Algorithm::CN_DOUBLE:
return CN_ITER * 2;

# ifdef XMRIG_ALGO_CN_PICO
case Algorithm::CN_PICO_0:
case Algorithm::CN_PICO_TLO:
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

default:
break;
Expand All @@ -117,13 +133,23 @@ class CnAlgo

inline static uint32_t mask(Algorithm::Id algo)
{
# ifdef XMRIG_ALGO_CN_PICO
if (algo == Algorithm::CN_PICO_0) {
return 0x1FFF0;
}
# 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

return ((memory(algo) - 1) / 16) * 16;
}
Expand All @@ -133,32 +159,47 @@ class CnAlgo
switch (algo) {
case Algorithm::CN_0:
case Algorithm::CN_XAO:
# ifdef XMRIG_ALGO_CN_LITE
case Algorithm::CN_LITE_0:
# endif
# ifdef XMRIG_ALGO_CN_HEAVY
case Algorithm::CN_HEAVY_0:
case Algorithm::CN_HEAVY_XHV:
# endif
case Algorithm::CN_CCX:
return Algorithm::CN_0;

case Algorithm::CN_1:
case Algorithm::CN_FAST:
case Algorithm::CN_RTO:
# ifdef XMRIG_ALGO_CN_LITE
case Algorithm::CN_LITE_1:
# endif
# ifdef XMRIG_ALGO_CN_HEAVY
case Algorithm::CN_HEAVY_TUBE:
# endif
return Algorithm::CN_1;

# ifdef XMRIG_ALGO_CN_GPU
case Algorithm::CN_GPU:
return Algorithm::CN_GPU;
# endif

case Algorithm::CN_2:
case Algorithm::CN_R:
case Algorithm::CN_HALF:
case Algorithm::CN_RWZ:
case Algorithm::CN_ZLS:
case Algorithm::CN_DOUBLE:
# ifdef XMRIG_ALGO_CN_PICO
case Algorithm::CN_PICO_0:
case Algorithm::CN_PICO_TLO:
# endif
# ifdef XMRIG_ALGO_CN_FEMTO
case Algorithm::CN_UPX2:
# endif
return Algorithm::CN_2;

case Algorithm::CN_GPU:
return Algorithm::CN_GPU;

default:
break;
}
Expand All @@ -183,6 +224,7 @@ template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_FAST>::base() con
template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_RTO>::base() const { return Algorithm::CN_1; }
template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_LITE_1>::base() const { return Algorithm::CN_1; }
template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_HEAVY_TUBE>::base() const { return Algorithm::CN_1; }
template<> constexpr inline Algorithm::Id CnAlgo<Algorithm::CN_UPX2>::base() const { return Algorithm::CN_2; }


template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_FAST>::iterations() const { return CN_ITER / 2; }
Expand All @@ -200,6 +242,7 @@ template<> constexpr inline uint32_t CnAlgo<Algorithm::CN_ZLS>::iterations() con
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_UPX2>::iterations() const { return CN_ITER / 32; }


template<> constexpr inline size_t CnAlgo<Algorithm::CN_LITE_0>::memory() const { return CN_MEMORY / 2; }
Expand All @@ -209,10 +252,12 @@ template<> constexpr inline size_t CnAlgo<Algorithm::CN_HEAVY_TUBE>::memory() co
template<> constexpr inline size_t CnAlgo<Algorithm::CN_HEAVY_XHV>::memory() const { return CN_MEMORY * 2; }
template<> constexpr inline size_t CnAlgo<Algorithm::CN_PICO_0>::memory() const { return CN_MEMORY / 8; }
template<> constexpr inline size_t CnAlgo<Algorithm::CN_PICO_TLO>::memory() const { return CN_MEMORY / 8; }
template<> constexpr inline size_t CnAlgo<Algorithm::CN_UPX2>::memory() const { return CN_MEMORY / 16; }


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; }


} /* namespace xmrig_cuda */
Expand Down
7 changes: 5 additions & 2 deletions src/crypto/common/Algorithm.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,8 +6,8 @@
* 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-2020 SChernykh <https://github.com/SChernykh>
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <[email protected]>
* Copyright 2018-2021 SChernykh <https://github.com/SChernykh>
* Copyright 2016-2021 XMRig <https://github.com/xmrig>, <[email protected]>
*
* 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
Expand Down Expand Up @@ -85,6 +85,9 @@ static AlgoName const algorithm_names[] = {
{ "cn-pico", Algorithm::CN_PICO_0 },
{ "cn-pico/tlo", Algorithm::CN_PICO_TLO },
# endif
# ifdef XMRIG_ALGO_CN_FEMTO
{ "cn/upx2", Algorithm::CN_UPX2 },
# endif
# ifdef XMRIG_ALGO_ASTROBWT
{ "astrobwt", Algorithm::ASTROBWT_DERO },
# endif
Expand Down
14 changes: 11 additions & 3 deletions src/crypto/common/Algorithm.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,8 +6,8 @@
* 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-2020 SChernykh <https://github.com/SChernykh>
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <[email protected]>
* Copyright 2018-2021 SChernykh <https://github.com/SChernykh>
* Copyright 2016-2021 XMRig <https://github.com/xmrig>, <[email protected]>
*
* 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
Expand Down Expand Up @@ -59,6 +59,7 @@ class Algorithm
CN_PICO_TLO, // "cn-pico/tlo" CryptoNight-Pico (TLO)
CN_CCX, // "cn/ccx" Conceal (CCX)
CN_GPU, // "cn/gpu" CryptoNight-GPU (Ryo).
CN_UPX2, // "cn/upx2" Uplexa (UPX2)
RX_0, // "rx/0" RandomX (reference configuration).
RX_WOW, // "rx/wow" RandomWOW (Wownero).
RX_ARQ, // "rx/arq" RandomARQ (Arqma).
Expand All @@ -79,6 +80,7 @@ class Algorithm
CN_LITE,
CN_HEAVY,
CN_PICO,
CN_FEMTO,
RANDOM_X,
ARGON2,
ASTROBWT,
Expand All @@ -90,7 +92,7 @@ class Algorithm
inline Algorithm(Id id) : m_id(id) {}
inline Algorithm(int id) : m_id(id > INVALID && id < MAX ? static_cast<Id>(id) : INVALID) {}

inline bool isCN() const { auto f = family(); return f == CN || f == CN_LITE || f == CN_HEAVY || f == CN_PICO; }
inline bool isCN() const { auto f = family(); return f == CN || f == CN_LITE || f == CN_HEAVY || f == CN_PICO || f == CN_FEMTO; }
inline bool isEqual(const Algorithm &other) const { return m_id == other.m_id; }
inline bool isValid() const { return m_id != INVALID; }
inline Family family() const { return family(m_id); }
Expand Down Expand Up @@ -144,6 +146,9 @@ class Algorithm
case CN_PICO:
return oneMiB / 4;

case CN_FEMTO:
return oneMiB / 8;

default:
break;
}
Expand Down Expand Up @@ -234,6 +239,9 @@ class Algorithm
case CN_PICO_TLO:
return CN_PICO;

case CN_UPX2:
return CN_FEMTO;

case RX_0:
case RX_WOW:
case RX_ARQ:
Expand Down
18 changes: 14 additions & 4 deletions src/cuda_core.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -351,8 +351,8 @@ __global__ void cryptonight_core_gpu_phase2_double(
__syncthreads();
# endif

myChunks[idx1 ^ 2 + sub] = ((ALGO == Algorithm::CN_RWZ) ? chunk1 : chunk3) + bx1;
myChunks[idx1 ^ 4 + sub] = ((ALGO == Algorithm::CN_RWZ) ? chunk3 : chunk1) + bx0;
myChunks[idx1 ^ 2 + sub] = (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? chunk1 : chunk3) + bx1;
myChunks[idx1 ^ 4 + sub] = (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? chunk3 : chunk1) + bx0;
myChunks[idx1 ^ 6 + sub] = chunk2 + ax0;
}

Expand Down Expand Up @@ -406,8 +406,8 @@ __global__ void cryptonight_core_gpu_phase2_double(
__syncthreads( );
# endif

myChunks[idx1 ^ 2 + sub] = ((ALGO == Algorithm::CN_RWZ) ? chunk1 : chunk3) + bx1;
myChunks[idx1 ^ 4 + sub] = ((ALGO == Algorithm::CN_RWZ) ? chunk3 : chunk1) + bx0;
myChunks[idx1 ^ 2 + sub] = (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? chunk1 : chunk3) + bx1;
myChunks[idx1 ^ 4 + sub] = (((ALGO == Algorithm::CN_RWZ) || (ALGO == Algorithm::CN_UPX2)) ? chunk3 : chunk1) + bx0;
myChunks[idx1 ^ 6 + sub] = chunk2 + ax0;

ax0 += res;
Expand Down Expand Up @@ -1020,4 +1020,14 @@ void cryptonight_gpu_hash(nvid_ctx *ctx, const xmrig_cuda::Algorithm &algorithm,
break;
}
}
else if (algorithm.family() == Algorithm::CN_FEMTO) {
switch (algorithm.id()) {
case Algorithm::CN_UPX2:
cryptonight_core_gpu_hash<Algorithm::CN_UPX2>(ctx, startNonce);
break;

default:
break;
}
}
}
8 changes: 4 additions & 4 deletions src/version.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -5,8 +5,8 @@
* 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-2020 SChernykh <https://github.com/SChernykh>
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <[email protected]>
* Copyright 2018-2021 SChernykh <https://github.com/SChernykh>
* Copyright 2016-2021 XMRig <https://github.com/xmrig>, <[email protected]>
*
* 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
Expand All @@ -28,13 +28,13 @@
#define APP_ID "xmrig-cuda"
#define APP_NAME "XMRig"
#define APP_DESC "XMRig CUDA plugin"
#define APP_VERSION "6.5.0-mo1"
#define APP_VERSION "6.12.0-mo1"
#define APP_DOMAIN "xmrig.com"
#define APP_SITE "www.xmrig.com"
#define APP_COPYRIGHT "Copyright (C) 2016-2020 xmrig.com"

#define APP_VER_MAJOR 6
#define APP_VER_MINOR 5
#define APP_VER_MINOR 12
#define APP_VER_PATCH 0

#define API_VERSION 3
Expand Down

1 comment on commit 58a437c

@protomens
Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Hello, I couldn't find the correct commit where this code lies, but here is my issue anyway.

I encounter the following error when trying to run xmrig-mo with the xmrig-mo-cuda. I was able to compile the cuda library just fine.

I'm on Ubuntu 20.04 with a NVIDIA GeForce GT 710

futex(0x55818eb18b0c, FUTEX_WAKE_PRIVATE, 1[2021-05-31 02:52:38.505] nvidia thread #0 failed with error <cryptonight_core_gpu_hash_gpu>:859 "invalid configuration argument"

I tried to isolate the problem, but couldn't quite figure it out without really diving in.

strace just loops through infinitely with the following:

epoll_wait(17, [], 1024, 500)           = 0
epoll_wait(17, [], 1024, 463)           = 0
epoll_wait(17, [], 1024, 36)            = 0
epoll_wait(17, [], 1024, 500)           = 0
epoll_wait(17, [], 1024, 464)           = 0
epoll_wait(17, [], 1024, 35)            = 0
epoll_wait(17, [], 1024, 500)           = 0
epoll_wait(17, [], 1024, 464)           = 0
epoll_wait(17, [], 1024, 36)            = 0
epoll_wait(17, [], 1024, 500)           = 0
epoll_wait(17, [], 1024, 463)           = 0
epoll_wait(17, [], 1024, 37)            = 0
epoll_wait(17, [], 1024, 500)           = 0
epoll_wait(17, [], 1024, 462)           = 0
epoll_wait(17, [], 1024, 37)            = 0
epoll_wait(17, [], 1024, 500)           = 0
epoll_wait(17, [], 1024, 463)           = 0
epoll_wait(17, [], 1024, 36)            = 0
epoll_wait(17, [], 1024, 500)           = 0

etc. Any help is appreciated. It's not a great video card, but every hash counts imho. Thanks!

Please sign in to comment.