diff options
30 files changed, 477 insertions, 134 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt index 1cc2ef6..595631d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -36,6 +36,19 @@ if(NOT CMAKE_BUILD_TYPE) endif() set_property(CACHE CMAKE_BUILD_TYPE PROPERTY STRINGS "${BUILD_TYPE}") +set(XMR-STAK_CURRENCY "all" CACHE STRING "select miner currency") +set_property(CACHE XMR-STAK_CURRENCY PROPERTY STRINGS "all;monero;aeon") + +if("${XMR-STAK_CURRENCY}" STREQUAL "all") + message(STATUS "Set miner currency to 'monero' and 'aeon'") +elseif("${XMR-STAK_CURRENCY}" STREQUAL "aeon") + message(STATUS "Set miner currency to 'aeon'") + add_definitions("-DCONF_NO_MONERO=1") +elseif("${XMR-STAK_CURRENCY}" STREQUAL "monero") + message(STATUS "Set miner currency to 'monero'") + add_definitions("-DCONF_NO_AEON=1") +endif() + # option to add static libgcc and libstdc++ option(CMAKE_LINK_STATIC "link as much as possible libraries static" OFF) @@ -1,6 +1,6 @@ -# XMR-Stak - Monero All-in-One Mining Software +# XMR-Stak - Monero/Aeon All-in-One Mining Software -XMR-Stak is a universal Stratum pool miner. This miner supports CPUs, AMD and NVIDIA gpus. +XMR-Stak is a universal Stratum pool miner. This miner supports CPUs, AMD and NVIDIA gpus and can be used to mine the crypto currency Monero and Aeon. ## HTML reports <img src="https://gist.githubusercontent.com/fireice-uk/2da301131ac01695ff79539a27b81d68/raw/4c09cdeee86f94df2e9dd86b927e64aded6184f5/xmr-stak-cpu-hashrate.png" width="260"> <img src="https://gist.githubusercontent.com/fireice-uk/2da301131ac01695ff79539a27b81d68/raw/4c09cdeee86f94df2e9dd86b927e64aded6184f5/xmr-stak-cpu-results.png" width="260"> <img src="https://gist.githubusercontent.com/fireice-uk/2da301131ac01695ff79539a27b81d68/raw/4c09cdeee86f94df2e9dd86b927e64aded6184f5/xmr-stak-cpu-connection.png" width="260"> @@ -18,6 +18,7 @@ XMR-Stak is a universal Stratum pool miner. This miner supports CPUs, AMD and NV - support all common backends (CPU/x86, AMD-GPU and NVIDIA-GPU) - support all common OS (Linux, Windows and MacOS) +- supports algorithm cryptonight for Monero (XMR) and cryptonight-light (AEON) - easy to use - guided start (no need to edit a config file for the first start) - auto configuration for each backend @@ -7,6 +7,7 @@ * [Error: MEMORY ALLOC FAILED: mmap failed](#error-memory-alloc-failed-mmap-failed) * [Illegal instruction (core dumped)](#illegal-instruction) * [Virus Protection Alert](#virus-protection-alert) +* [Change Currency to Mine](#change-currency-to-mine) ## SeLockMemoryPrivilege failed @@ -59,3 +60,9 @@ Some Virus protection software flag the miner binary as *Male Ware*. In this case the binary is moved to the quarantine area of the protection software. This is a wrong alert and not avoid by use. Add the binary to to protection software white list to solve this issue.s + +## Change Currency to Mine + +If the miner is compiled for Monero and Aeon than you can change + - the value `currency` in the config *or* + - start the miner with the [command line option](usage.md) `--currency monero` or `--currency aeon` diff --git a/doc/compile.md b/doc/compile.md index 248d6d9..537a736 100644 --- a/doc/compile.md +++ b/doc/compile.md @@ -47,6 +47,7 @@ After the configuration you need to compile the miner, follow the guide for your - `WIN_UAC` will enable or disable the "Run As Administrator" prompt on Windows. - UAC confirmation is needed to use large pages on Windows 7. - On Windows 10 it is only needed once to set up the account to use them. +- `XMR-STAK_CURRENCY` - compile for Monero(XMR) or Aeon(AEON) usage only e.g. `cmake .. -DXMR-STAK_CURRENCY=monero` ## CPU Build Options diff --git a/doc/usage.md b/doc/usage.md index 42a29fd..9a38384 100644 --- a/doc/usage.md +++ b/doc/usage.md @@ -31,8 +31,9 @@ The miner allow to overwrite some of the settings via command line options. ``` Usage: xmr-stak [OPTION]... - -c, --config FILE common miner configuration file -h, --help show this help + -c, --config FILE common miner configuration file + --currency NAME currency to mine: monero or aeon --noCPU disable the CPU miner backend --cpu FILE CPU backend miner config file --noAMD disable the AMD miner backend diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index 3575854..22ce5d0 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -13,6 +13,9 @@ * along with this program. If not, see <http://www.gnu.org/licenses/>. */ +#include "xmrstak/backend/cryptonight.hpp" +#include "xmrstak/jconf.hpp" + #include <stdio.h> #include <string.h> #include <math.h> @@ -245,8 +248,24 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_ return ERR_OCL_API; } + size_t hashMemSize; + int threadMemMask; + int hasIterations; + if(::jconf::inst()->IsCurrencyMonero()) + { + hashMemSize = MONERO_MEMORY; + threadMemMask = MONERO_MASK; + hasIterations = MONERO_ITER; + } + else + { + hashMemSize = AEON_MEMORY; + threadMemMask = AEON_MASK; + hasIterations = AEON_ITER; + } + size_t g_thd = ctx->rawIntensity; - ctx->ExtraBuffers[0] = clCreateBuffer(opencl_ctx, CL_MEM_READ_WRITE, (1 << 21) * g_thd, NULL, &ret); + ctx->ExtraBuffers[0] = clCreateBuffer(opencl_ctx, CL_MEM_READ_WRITE, hashMemSize * g_thd, NULL, &ret); if(ret != CL_SUCCESS) { printer::inst()->print_msg(L1,"Error %s when calling clCreateBuffer to create hash scratchpads buffer.", err_to_str(ret)); @@ -307,8 +326,9 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_ return ERR_OCL_API; } - char options[32]; - snprintf(options, sizeof(options), "-I. -DWORKSIZE=%llu", int_port(ctx->workSize)); + char options[256]; + snprintf(options, sizeof(options), + "-DITERATIONS=%d -DMASK=%d -DWORKSIZE=%llu", hasIterations, threadMemMask, int_port(ctx->workSize)); ret = clBuildProgram(ctx->Program, 1, &ctx->DeviceID, options, NULL, NULL); if(ret != CL_SUCCESS) { diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl index a6a5910..966199b 100644 --- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl +++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl @@ -91,7 +91,7 @@ static const __constant ulong keccakf_rndc[24] = 0x8000000080008081, 0x8000000000008009, 0x000000000000008a, 0x0000000000000088, 0x0000000080008009, 0x000000008000000a, 0x000000008000808b, 0x800000000000008b, 0x8000000000008089, - 0x8000000000008003, 0x8000000000008002, 0x8000000000000080, + 0x8000000000008003, 0x8000000000008002, 0x8000000000000080, 0x000000000000800a, 0x800000008000000a, 0x8000000080008081, 0x8000000000008080, 0x0000000080000001, 0x8000000080008008 }; @@ -440,7 +440,7 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul if(gIdx < Threads) { states += 25 * gIdx; - Scratchpad += gIdx * (0x80000 >> 2); + Scratchpad += gIdx * (ITERATIONS >> 2); ((ulong8 *)State)[0] = vload8(0, input); State[8] = input[8]; @@ -482,7 +482,7 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul if(gIdx < Threads) { #pragma unroll 2 - for(int i = 0; i < 0x4000; ++i) + for(int i = 0; i < (ITERATIONS >> 5); ++i) { #pragma unroll for(int j = 0; j < 10; ++j) @@ -519,7 +519,7 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre if(gIdx < Threads) { states += 25 * gIdx; - Scratchpad += gIdx * (0x80000 >> 2); + Scratchpad += gIdx * (ITERATIONS >> 2); a[0] = states[0] ^ states[4]; b[0] = states[2] ^ states[6]; @@ -535,23 +535,23 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre if(gIdx < Threads) { #pragma unroll 8 - for(int i = 0; i < 0x80000; ++i) + for(int i = 0; i < ITERATIONS; ++i) { ulong c[2]; - ((uint4 *)c)[0] = Scratchpad[IDX((a[0] & 0x1FFFF0) >> 4)]; + ((uint4 *)c)[0] = Scratchpad[IDX((a[0] & MASK) >> 4)]; ((uint4 *)c)[0] = AES_Round(AES0, AES1, AES2, AES3, ((uint4 *)c)[0], ((uint4 *)a)[0]); //b_x ^= ((uint4 *)c)[0]; - Scratchpad[IDX((a[0] & 0x1FFFF0) >> 4)] = b_x ^ ((uint4 *)c)[0]; + Scratchpad[IDX((a[0] & MASK) >> 4)] = b_x ^ ((uint4 *)c)[0]; uint4 tmp; - tmp = Scratchpad[IDX((c[0] & 0x1FFFF0) >> 4)]; + tmp = Scratchpad[IDX((c[0] & MASK) >> 4)]; a[1] += c[0] * as_ulong2(tmp).s0; a[0] += mul_hi(c[0], as_ulong2(tmp).s0); - Scratchpad[IDX((c[0] & 0x1FFFF0) >> 4)] = ((uint4 *)a)[0]; + Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0]; ((uint4 *)a)[0] ^= tmp; @@ -588,7 +588,7 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u if(gIdx < Threads) { states += 25 * gIdx; - Scratchpad += gIdx * (0x80000 >> 2); + Scratchpad += gIdx * (ITERATIONS >> 2); #if defined(__Tahiti__) || defined(__Pitcairn__) @@ -611,7 +611,7 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u if(gIdx < Threads) { #pragma unroll 2 - for(int i = 0; i < 0x4000; ++i) + for(int i = 0; i < (ITERATIONS >> 5); ++i) { text ^= Scratchpad[IDX((i << 3) + get_local_id(1))]; diff --git a/xmrstak/backend/amd/autoAdjust.hpp b/xmrstak/backend/amd/autoAdjust.hpp index 2a22a08..87e6299 100644 --- a/xmrstak/backend/amd/autoAdjust.hpp +++ b/xmrstak/backend/amd/autoAdjust.hpp @@ -8,6 +8,8 @@ #include "xmrstak/misc/console.hpp" #include "xmrstak/misc/configEditor.hpp" #include "xmrstak/params.hpp" +#include "xmrstak/backend/cryptonight.hpp" +#include "xmrstak/jconf.hpp" #include <vector> #include <cstdio> @@ -81,6 +83,16 @@ private: constexpr size_t byteToMiB = 1024u * 1024u; + size_t hashMemSize; + if(::jconf::inst()->IsCurrencyMonero()) + { + hashMemSize = MONERO_MEMORY; + } + else + { + hashMemSize = AEON_MEMORY; + } + std::string conf; int i = 0; for(auto& ctx : devVec) @@ -88,7 +100,7 @@ private: // keep 64MiB memory free (value is randomly chosen) size_t availableMem = ctx.freeMem - (64u * 1024 * 1024); // 224byte extra memory is used per thread for meta data - size_t perThread = (size_t(1u)<<21) + 224u; + size_t perThread = hashMemSize + 224u; size_t max_intensity = availableMem / perThread; // 1000 is a magic selected limit \todo select max intensity depending of the gpu type size_t possibleIntensity = std::min( size_t(1000u) , max_intensity ); diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp index e262e0e..f12e12c 100644 --- a/xmrstak/backend/amd/minethd.cpp +++ b/xmrstak/backend/amd/minethd.cpp @@ -183,7 +183,7 @@ void minethd::work_main() uint64_t iCount = 0; cryptonight_ctx* cpu_ctx; cpu_ctx = cpu::minethd::minethd_alloc_ctx(); - cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/); + cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, ::jconf::inst()->IsCurrencyMonero()); globalStates::inst().iConsumeCnt++; while (bQuit == 0) diff --git a/xmrstak/backend/cpu/autoAdjust.hpp b/xmrstak/backend/cpu/autoAdjust.hpp index 294bc6f..7bdb14e 100644 --- a/xmrstak/backend/cpu/autoAdjust.hpp +++ b/xmrstak/backend/cpu/autoAdjust.hpp @@ -6,6 +6,7 @@ #include "xmrstak/jconf.hpp" #include "xmrstak/misc/configEditor.hpp" #include "xmrstak/params.hpp" +#include "xmrstak/backend/cryptonight.hpp" #include <string> #ifdef _WIN32 @@ -32,8 +33,21 @@ class autoAdjust { public: + size_t hashMemSize; + size_t halfHashMemSize; + autoAdjust() { + if(::jconf::inst()->IsCurrencyMonero()) + { + hashMemSize = MONERO_MEMORY; + halfHashMemSize = hashMemSize / 2u; + } + else + { + hashMemSize = AEON_MEMORY; + halfHashMemSize = hashMemSize / 2u; + } } bool printConfig() @@ -49,9 +63,9 @@ public: std::string conf; - if(!detectL3Size() || L3KB_size < 1024 || L3KB_size > 102400) + if(!detectL3Size() || L3KB_size < halfHashMemSize || L3KB_size > (halfHashMemSize * 100u)) { - if(L3KB_size < 1024 || L3KB_size > 102400) + if(L3KB_size < halfHashMemSize || L3KB_size > (halfHashMemSize * 100)) printer::inst()->print_msg(L0, "Autoconf failed: L3 size sanity check failed - %u KB.", L3KB_size); conf += std::string(" { \"low_power_mode\" : false, \"no_prefetch\" : true, \"affine_to_cpu\" : false },\n"); @@ -74,7 +88,7 @@ public: if(L3KB_size <= 0) break; - double_mode = L3KB_size / 2048 > (int32_t)(corecnt-i); + double_mode = L3KB_size / hashMemSize > (int32_t)(corecnt-i); conf += std::string(" { \"low_power_mode\" : "); conf += std::string(double_mode ? "true" : "false"); @@ -93,9 +107,9 @@ public: aff_id++; if(double_mode) - L3KB_size -= 4096; + L3KB_size -= hashMemSize * 2u; else - L3KB_size -= 2048; + L3KB_size -= hashMemSize; } } @@ -128,7 +142,7 @@ private: } L3KB_size = ((get_masked(cpu_info[1], 31, 22) + 1) * (get_masked(cpu_info[1], 21, 12) + 1) * - (get_masked(cpu_info[1], 11, 0) + 1) * (cpu_info[2] + 1)) / 1024; + (get_masked(cpu_info[1], 11, 0) + 1) * (cpu_info[2] + 1)) / halfHashMemSize; return true; } diff --git a/xmrstak/backend/cpu/autoAdjustHwloc.hpp b/xmrstak/backend/cpu/autoAdjustHwloc.hpp index ad3b863..ddeb89b 100644 --- a/xmrstak/backend/cpu/autoAdjustHwloc.hpp +++ b/xmrstak/backend/cpu/autoAdjustHwloc.hpp @@ -3,6 +3,7 @@ #include "xmrstak/misc/console.hpp" #include "xmrstak/misc/configEditor.hpp" #include "xmrstak/params.hpp" +#include "xmrstak/backend/cryptonight.hpp" #ifdef _WIN32 #include <windows.h> @@ -27,6 +28,16 @@ public: autoAdjust() { + if(::jconf::inst()->IsCurrencyMonero()) + { + hashMemSize = MONERO_MEMORY; + halfHashMemSize = hashMemSize / 2u; + } + else + { + hashMemSize = AEON_MEMORY; + halfHashMemSize = hashMemSize / 2u; + } } bool printConfig() @@ -86,7 +97,9 @@ public: } private: - static constexpr size_t hashSize = 2 * 1024 * 1024; + size_t hashMemSize; + size_t halfHashMemSize; + std::vector<uint32_t> results; template<typename func> @@ -161,8 +174,8 @@ private: { hwloc_obj_t l2obj = obj->children[i]; //If L2 is exclusive and greater or equal to 2MB add room for one more hash - if(isCacheObject(l2obj) && l2obj->attr != nullptr && l2obj->attr->cache.size >= hashSize) - cacheSize += hashSize; + if(isCacheObject(l2obj) && l2obj->attr != nullptr && l2obj->attr->cache.size >= hashMemSize) + cacheSize += hashMemSize; } } @@ -170,7 +183,7 @@ private: cores.reserve(16); findChildrenByType(obj, HWLOC_OBJ_CORE, [&cores](hwloc_obj_t found) { cores.emplace_back(found); } ); - size_t cacheHashes = (cacheSize + hashSize/2) / hashSize; + size_t cacheHashes = (cacheSize + halfHashMemSize) / hashMemSize; //Firstly allocate PU 0 of every CORE, then PU 1 etc. size_t pu_id = 0; diff --git a/xmrstak/backend/cpu/crypto/cryptonight.h b/xmrstak/backend/cpu/crypto/cryptonight.h index 978c798..631c39a 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight.h +++ b/xmrstak/backend/cpu/crypto/cryptonight.h @@ -7,8 +7,8 @@ extern "C" { #include <stddef.h> #include <inttypes.h> +#include "xmrstak/backend/cryptonight.hpp" -#define MEMORY 2097152 typedef struct { uint8_t hash_state[224]; // Need only 200, explicit align diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h index 942d511..2a6a769 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h +++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h @@ -287,7 +287,7 @@ void cn_implode_scratchpad(const __m128i* input, __m128i* output) _mm_store_si128(output + 11, xout7); } -template<size_t ITERATIONS, size_t MEM, bool SOFT_AES, bool PREFETCH> +template<size_t MASK, size_t ITERATIONS, size_t MEM, bool SOFT_AES, bool PREFETCH> void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_ctx* ctx0) { keccak((const uint8_t *)input, len, ctx0->hash_state, 200); @@ -308,36 +308,36 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c for(size_t i = 0; i < ITERATIONS; i++) { __m128i cx; - cx = _mm_load_si128((__m128i *)&l0[idx0 & 0x1FFFF0]); + cx = _mm_load_si128((__m128i *)&l0[idx0 & MASK]); if(SOFT_AES) cx = soft_aesenc(cx, _mm_set_epi64x(ah0, al0)); else cx = _mm_aesenc_si128(cx, _mm_set_epi64x(ah0, al0)); - _mm_store_si128((__m128i *)&l0[idx0 & 0x1FFFF0], _mm_xor_si128(bx0, cx)); + _mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx)); idx0 = _mm_cvtsi128_si64(cx); bx0 = cx; if(PREFETCH) - _mm_prefetch((const char*)&l0[idx0 & 0x1FFFF0], _MM_HINT_T0); + _mm_prefetch((const char*)&l0[idx0 & MASK], _MM_HINT_T0); uint64_t hi, lo, cl, ch; - cl = ((uint64_t*)&l0[idx0 & 0x1FFFF0])[0]; - ch = ((uint64_t*)&l0[idx0 & 0x1FFFF0])[1]; + cl = ((uint64_t*)&l0[idx0 & MASK])[0]; + ch = ((uint64_t*)&l0[idx0 & MASK])[1]; lo = _umul128(idx0, cl, &hi); al0 += hi; ah0 += lo; - ((uint64_t*)&l0[idx0 & 0x1FFFF0])[0] = al0; - ((uint64_t*)&l0[idx0 & 0x1FFFF0])[1] = ah0; + ((uint64_t*)&l0[idx0 & MASK])[0] = al0; + ((uint64_t*)&l0[idx0 & MASK])[1] = ah0; ah0 ^= ch; al0 ^= cl; idx0 = al0; if(PREFETCH) - _mm_prefetch((const char*)&l0[idx0 & 0x1FFFF0], _MM_HINT_T0); + _mm_prefetch((const char*)&l0[idx0 & MASK], _MM_HINT_T0); } // Optim - 90% time boundary @@ -352,7 +352,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c // This lovely creation will do 2 cn hashes at a time. We have plenty of space on silicon // to fit temporary vars for two contexts. Function will read len*2 from input and write 64 bytes to output // We are still limited by L3 cache, so doubling will only work with CPUs where we have more than 2MB to core (Xeons) -template<size_t ITERATIONS, size_t MEM, bool SOFT_AES, bool PREFETCH> +template<size_t MASK, size_t ITERATIONS, size_t MEM, bool SOFT_AES, bool PREFETCH> void cryptonight_double_hash(const void* input, size_t len, void* output, cryptonight_ctx* __restrict ctx0, cryptonight_ctx* __restrict ctx1) { keccak((const uint8_t *)input, len, ctx0->hash_state, 200); @@ -381,66 +381,66 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto for (size_t i = 0; i < ITERATIONS; i++) { __m128i cx; - cx = _mm_load_si128((__m128i *)&l0[idx0 & 0x1FFFF0]); + cx = _mm_load_si128((__m128i *)&l0[idx0 & MASK]); if(SOFT_AES) cx = soft_aesenc(cx, _mm_set_epi64x(axh0, axl0)); else cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh0, axl0)); - _mm_store_si128((__m128i *)&l0[idx0 & 0x1FFFF0], _mm_xor_si128(bx0, cx)); + _mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx)); idx0 = _mm_cvtsi128_si64(cx); bx0 = cx; if(PREFETCH) - _mm_prefetch((const char*)&l0[idx0 & 0x1FFFF0], _MM_HINT_T0); + _mm_prefetch((const char*)&l0[idx0 & MASK], _MM_HINT_T0); - cx = _mm_load_si128((__m128i *)&l1[idx1 & 0x1FFFF0]); + cx = _mm_load_si128((__m128i *)&l1[idx1 & MASK]); if(SOFT_AES) cx = soft_aesenc(cx, _mm_set_epi64x(axh1, axl1)); else cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh1, axl1)); - _mm_store_si128((__m128i *)&l1[idx1 & 0x1FFFF0], _mm_xor_si128(bx1, cx)); + _mm_store_si128((__m128i *)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx)); idx1 = _mm_cvtsi128_si64(cx); bx1 = cx; if(PREFETCH) - _mm_prefetch((const char*)&l1[idx1 & 0x1FFFF0], _MM_HINT_T0); + _mm_prefetch((const char*)&l1[idx1 & MASK], _MM_HINT_T0); uint64_t hi, lo, cl, ch; - cl = ((uint64_t*)&l0[idx0 & 0x1FFFF0])[0]; - ch = ((uint64_t*)&l0[idx0 & 0x1FFFF0])[1]; + cl = ((uint64_t*)&l0[idx0 & MASK])[0]; + ch = ((uint64_t*)&l0[idx0 & MASK])[1]; lo = _umul128(idx0, cl, &hi); axl0 += hi; axh0 += lo; - ((uint64_t*)&l0[idx0 & 0x1FFFF0])[0] = axl0; - ((uint64_t*)&l0[idx0 & 0x1FFFF0])[1] = axh0; + ((uint64_t*)&l0[idx0 & MASK])[0] = axl0; + ((uint64_t*)&l0[idx0 & MASK])[1] = axh0; axh0 ^= ch; axl0 ^= cl; idx0 = axl0; if(PREFETCH) - _mm_prefetch((const char*)&l0[idx0 & 0x1FFFF0], _MM_HINT_T0); + _mm_prefetch((const char*)&l0[idx0 & MASK], _MM_HINT_T0); - cl = ((uint64_t*)&l1[idx1 & 0x1FFFF0])[0]; - ch = ((uint64_t*)&l1[idx1 & 0x1FFFF0])[1]; + cl = ((uint64_t*)&l1[idx1 & MASK])[0]; + ch = ((uint64_t*)&l1[idx1 & MASK])[1]; lo = _umul128(idx1, cl, &hi); axl1 += hi; axh1 += lo; - ((uint64_t*)&l1[idx1 & 0x1FFFF0])[0] = axl1; - ((uint64_t*)&l1[idx1 & 0x1FFFF0])[1] = axh1; + ((uint64_t*)&l1[idx1 & MASK])[0] = axl1; + ((uint64_t*)&l1[idx1 & MASK])[1] = axh1; axh1 ^= ch; axl1 ^= cl; idx1 = axl1; if(PREFETCH) - _mm_prefetch((const char*)&l1[idx1 & 0x1FFFF0], _MM_HINT_T0); + _mm_prefetch((const char*)&l1[idx1 & MASK], _MM_HINT_T0); } // Optim - 90% time boundary diff --git a/xmrstak/backend/cpu/crypto/cryptonight_common.cpp b/xmrstak/backend/cpu/crypto/cryptonight_common.cpp index 0690415..c73dbd8 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight_common.cpp +++ b/xmrstak/backend/cpu/crypto/cryptonight_common.cpp @@ -30,6 +30,8 @@ extern "C" } #include "cryptonight.h" #include "cryptonight_aesni.h" +#include "xmrstak/backend/cryptonight.hpp" +#include "xmrstak/jconf.hpp" #include <stdio.h> #include <stdlib.h> @@ -194,12 +196,21 @@ size_t cryptonight_init(size_t use_fast_mem, size_t use_mlock, alloc_msg* msg) cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, alloc_msg* msg) { + size_t hashMemSize; + if(::jconf::inst()->IsCurrencyMonero()) + { + hashMemSize = MONERO_MEMORY; + } + else + { + hashMemSize = AEON_MEMORY; + } cryptonight_ctx* ptr = (cryptonight_ctx*)_mm_malloc(sizeof(cryptonight_ctx), 4096); if(use_fast_mem == 0) { // use 2MiB aligned memory - ptr->long_state = (uint8_t*)_mm_malloc(MEMORY, 2*1024*1024); + ptr->long_state = (uint8_t*)_mm_malloc(hashMemSize, hashMemSize); ptr->ctx_info[0] = 0; ptr->ctx_info[1] = 0; return ptr; @@ -208,7 +219,7 @@ cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, al #ifdef _WIN32 SIZE_T iLargePageMin = GetLargePageMinimum(); - if(MEMORY > iLargePageMin) + if(hashMemSize > iLargePageMin) iLargePageMin *= 2; ptr->long_state = (uint8_t*)VirtualAlloc(NULL, iLargePageMin, @@ -231,13 +242,13 @@ cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, al #else #if defined(__APPLE__) - ptr->long_state = (uint8_t*)mmap(0, MEMORY, PROT_READ | PROT_WRITE, + ptr->long_state = (uint8_t*)mmap(0, hashMemSize, PROT_READ | PROT_WRITE, MAP_PRIVATE | MAP_ANON, VM_FLAGS_SUPERPAGE_SIZE_2MB, 0); #elif defined(__FreeBSD__) - ptr->long_state = (uint8_t*)mmap(0, MEMORY, PROT_READ | PROT_WRITE, + ptr->long_state = (uint8_t*)mmap(0, hashMemSize, PROT_READ | PROT_WRITE, MAP_PRIVATE | MAP_ANONYMOUS | MAP_ALIGNED_SUPER | MAP_PREFAULT_READ, -1, 0); #else - ptr->long_state = (uint8_t*)mmap(0, MEMORY, PROT_READ | PROT_WRITE, + ptr->long_state = (uint8_t*)mmap(0, hashMemSize, PROT_READ | PROT_WRITE, MAP_PRIVATE | MAP_ANONYMOUS | MAP_HUGETLB | MAP_POPULATE, 0, 0); #endif @@ -250,11 +261,11 @@ cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, al ptr->ctx_info[0] = 1; - if(madvise(ptr->long_state, MEMORY, MADV_RANDOM|MADV_WILLNEED) != 0) + if(madvise(ptr->long_state, hashMemSize, MADV_RANDOM|MADV_WILLNEED) != 0) msg->warning = "madvise failed"; ptr->ctx_info[1] = 0; - if(use_mlock != 0 && mlock(ptr->long_state, MEMORY) != 0) + if(use_mlock != 0 && mlock(ptr->long_state, hashMemSize) != 0) msg->warning = "mlock failed"; else ptr->ctx_info[1] = 1; @@ -265,14 +276,23 @@ cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, al void cryptonight_free_ctx(cryptonight_ctx* ctx) { + size_t hashMemSize; + if(::jconf::inst()->IsCurrencyMonero()) + { + hashMemSize = MONERO_MEMORY; + } + else + { + hashMemSize = AEON_MEMORY; + } if(ctx->ctx_info[0] != 0) { #ifdef _WIN32 VirtualFree(ctx->long_state, 0, MEM_RELEASE); #else if(ctx->ctx_info[1] != 0) - munlock(ctx->long_state, MEMORY); - munmap(ctx->long_state, MEMORY); + munlock(ctx->long_state, hashMemSize); + munmap(ctx->long_state, hashMemSize); #endif // _WIN32 } else diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp index 1524eca..d36ebf1 100644 --- a/xmrstak/backend/cpu/minethd.cpp +++ b/xmrstak/backend/cpu/minethd.cpp @@ -199,30 +199,34 @@ bool minethd::self_test() return false; } - unsigned char out[64]; - bool bResult; + bool bResult = true; - cn_hash_fun hashf; - cn_hash_fun_dbl hashdf; + bool mineMonero = ::jconf::inst()->IsCurrencyMonero(); + if(mineMonero) + { + unsigned char out[64]; + cn_hash_fun hashf; + cn_hash_fun_dbl hashdf; - hashf = func_selector(::jconf::inst()->HaveHardwareAes(), false); - hashf("This is a test", 14, out, ctx0); - bResult = memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 32) == 0; - hashf = func_selector(::jconf::inst()->HaveHardwareAes(), true); - hashf("This is a test", 14, out, ctx0); - bResult &= memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 32) == 0; + hashf = func_selector(::jconf::inst()->HaveHardwareAes(), false, mineMonero); + hashf("This is a test", 14, out, ctx0); + bResult = memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 32) == 0; - hashdf = func_dbl_selector(::jconf::inst()->HaveHardwareAes(), false); - hashdf("The quick brown fox jumps over the lazy dogThe quick brown fox jumps over the lazy log", 43, out, ctx0, ctx1); - bResult &= memcmp(out, "\x3e\xbb\x7f\x9f\x7d\x27\x3d\x7c\x31\x8d\x86\x94\x77\x55\x0c\xc8\x00\xcf\xb1\x1b\x0c\xad\xb7\xff\xbd\xf6\xf8\x9f\x3a\x47\x1c\x59" - "\xb4\x77\xd5\x02\xe4\xd8\x48\x7f\x42\xdf\xe3\x8e\xed\x73\x81\x7a\xda\x91\xb7\xe2\x63\xd2\x91\x71\xb6\x5c\x44\x3a\x01\x2a\x41\x22", 64) == 0; + hashf = func_selector(::jconf::inst()->HaveHardwareAes(), true, mineMonero); + hashf("This is a test", 14, out, ctx0); + bResult &= memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 32) == 0; - hashdf = func_dbl_selector(::jconf::inst()->HaveHardwareAes(), true); - hashdf("The quick brown fox jumps over the lazy dogThe quick brown fox jumps over the lazy log", 43, out, ctx0, ctx1); - bResult &= memcmp(out, "\x3e\xbb\x7f\x9f\x7d\x27\x3d\x7c\x31\x8d\x86\x94\x77\x55\x0c\xc8\x00\xcf\xb1\x1b\x0c\xad\xb7\xff\xbd\xf6\xf8\x9f\x3a\x47\x1c\x59" - "\xb4\x77\xd5\x02\xe4\xd8\x48\x7f\x42\xdf\xe3\x8e\xed\x73\x81\x7a\xda\x91\xb7\xe2\x63\xd2\x91\x71\xb6\x5c\x44\x3a\x01\x2a\x41\x22", 64) == 0; + hashdf = func_dbl_selector(::jconf::inst()->HaveHardwareAes(), false, mineMonero); + hashdf("The quick brown fox jumps over the lazy dogThe quick brown fox jumps over the lazy log", 43, out, ctx0, ctx1); + bResult &= memcmp(out, "\x3e\xbb\x7f\x9f\x7d\x27\x3d\x7c\x31\x8d\x86\x94\x77\x55\x0c\xc8\x00\xcf\xb1\x1b\x0c\xad\xb7\xff\xbd\xf6\xf8\x9f\x3a\x47\x1c\x59" + "\xb4\x77\xd5\x02\xe4\xd8\x48\x7f\x42\xdf\xe3\x8e\xed\x73\x81\x7a\xda\x91\xb7\xe2\x63\xd2\x91\x71\xb6\x5c\x44\x3a\x01\x2a\x41\x22", 64) == 0; + hashdf = func_dbl_selector(::jconf::inst()->HaveHardwareAes(), true, mineMonero); + hashdf("The quick brown fox jumps over the lazy dogThe quick brown fox jumps over the lazy log", 43, out, ctx0, ctx1); + bResult &= memcmp(out, "\x3e\xbb\x7f\x9f\x7d\x27\x3d\x7c\x31\x8d\x86\x94\x77\x55\x0c\xc8\x00\xcf\xb1\x1b\x0c\xad\xb7\xff\xbd\xf6\xf8\x9f\x3a\x47\x1c\x59" + "\xb4\x77\xd5\x02\xe4\xd8\x48\x7f\x42\xdf\xe3\x8e\xed\x73\x81\x7a\xda\x91\xb7\xe2\x63\xd2\x91\x71\xb6\x5c\x44\x3a\x01\x2a\x41\x22", 64) == 0; + } cryptonight_free_ctx(ctx0); cryptonight_free_ctx(ctx1); @@ -285,24 +289,47 @@ void minethd::consume_work() globalStates::inst().inst().iConsumeCnt++; } -minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch) +minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, bool mineMonero) { // We have two independent flag bits in the functions // therefore we will build a binary digit and select the // function as a two digit binary - // Digit order SOFT_AES, NO_PREFETCH - - static const cn_hash_fun func_table[4] = { - cryptonight_hash<0x80000, MEMORY, false, false>, - cryptonight_hash<0x80000, MEMORY, false, true>, - cryptonight_hash<0x80000, MEMORY, true, false>, - cryptonight_hash<0x80000, MEMORY, true, true> + // Digit order SOFT_AES, NO_PREFETCH, MINER_ALGO + + static const cn_hash_fun func_table[] = { + /* there will be 8 function entries if `CONF_NO_MONERO` and `CONF_NO_AEON` + * is not defined. If one is defined there will be 4 entries. + */ +#ifndef CONF_NO_MONERO + cryptonight_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, false>, + cryptonight_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, true>, + cryptonight_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, false>, + cryptonight_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, true> +#endif +#if (!defined(CONF_NO_AEON)) && (!defined(CONF_NO_MONERO)) + // comma will be added only if Monero and Aeon is build + , +#endif +#ifndef CONF_NO_AEON + cryptonight_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, false>, + cryptonight_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, true>, + cryptonight_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, false>, + cryptonight_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, true> +#endif }; - std::bitset<2> digit; + std::bitset<3> digit; digit.set(0, !bNoPrefetch); digit.set(1, !bHaveAes); + // define aeon settings +#if defined(CONF_NO_AEON) || defined(CONF_NO_MONERO) + // ignore 3rd bit if only on currency is active + digit.set(2, 0); +#else + digit.set(2, !mineMonero); +#endif + return func_table[digit.to_ulong()]; } @@ -320,7 +347,7 @@ void minethd::work_main() uint32_t* piNonce; job_result result; - hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch); + hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->IsCurrencyMonero()); ctx = minethd_alloc_ctx(); piHashVal = (uint64_t*)(result.bResult + 24); @@ -382,24 +409,47 @@ void minethd::work_main() cryptonight_free_ctx(ctx); } -minethd::cn_hash_fun_dbl minethd::func_dbl_selector(bool bHaveAes, bool bNoPrefetch) +minethd::cn_hash_fun_dbl minethd::func_dbl_selector(bool bHaveAes, bool bNoPrefetch, bool mineMonero) { // We have two independent flag bits in the functions // therefore we will build a binary digit and select the // function as a two digit binary - // Digit order SOFT_AES, NO_PREFETCH - - static const cn_hash_fun_dbl func_table[4] = { - cryptonight_double_hash<0x80000, MEMORY, false, false>, - cryptonight_double_hash<0x80000, MEMORY, false, true>, - cryptonight_double_hash<0x80000, MEMORY, true, false>, - cryptonight_double_hash<0x80000, MEMORY, true, true> + // Digit order SOFT_AES, NO_PREFETCH, MINER_ALGO + + static const cn_hash_fun_dbl func_table[] = { + /* there will be 8 function entries if `CONF_NO_MONERO` and `CONF_NO_AEON` + * is not defined. If one is defined there will be 4 entries. + */ +#ifndef CONF_NO_MONERO + cryptonight_double_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, false>, + cryptonight_double_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, true>, + cryptonight_double_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, false>, + cryptonight_double_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, true> +#endif +#if (!defined(CONF_NO_AEON)) && (!defined(CONF_NO_MONERO)) + // comma will be added only if Monero and Aeon is build + , +#endif +#ifndef CONF_NO_AEON + cryptonight_double_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, false>, + cryptonight_double_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, true>, + cryptonight_double_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, false>, + cryptonight_double_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, true> +#endif }; - std::bitset<2> digit; + std::bitset<3> digit; digit.set(0, !bNoPrefetch); digit.set(1, !bHaveAes); + // define aeon settings +#if defined(CONF_NO_AEON) || defined(CONF_NO_MONERO) + // ignore 3rd bit if only on currency is active + digit.set(2, 0); +#else + digit.set(2, !mineMonero); +#endif + return func_table[digit.to_ulong()]; } @@ -428,7 +478,7 @@ void minethd::double_work_main() uint32_t iNonce; job_result res; - hash_fun = func_dbl_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch); + hash_fun = func_dbl_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->IsCurrencyMonero()); ctx0 = minethd_alloc_ctx(); ctx1 = minethd_alloc_ctx(); diff --git a/xmrstak/backend/cpu/minethd.hpp b/xmrstak/backend/cpu/minethd.hpp index 5ffd44e..670ec8d 100644 --- a/xmrstak/backend/cpu/minethd.hpp +++ b/xmrstak/backend/cpu/minethd.hpp @@ -23,7 +23,7 @@ public: typedef void (*cn_hash_fun)(const void*, size_t, void*, cryptonight_ctx*); - static cn_hash_fun func_selector(bool bHaveAes, bool bNoPrefetch); + static cn_hash_fun func_selector(bool bHaveAes, bool bNoPrefetch, bool mineMonero); static bool thd_setaffinity(std::thread::native_handle_type h, uint64_t cpu_id); static cryptonight_ctx* minethd_alloc_ctx(); @@ -31,7 +31,7 @@ public: private: typedef void (*cn_hash_fun_dbl)(const void*, size_t, void*, cryptonight_ctx* __restrict, cryptonight_ctx* __restrict); - static cn_hash_fun_dbl func_dbl_selector(bool bHaveAes, bool bNoPrefetch); + static cn_hash_fun_dbl func_dbl_selector(bool bHaveAes, bool bNoPrefetch, bool mineMonero); minethd(miner_work& pWork, size_t iNo, bool double_work, bool no_prefetch, int64_t affinity); diff --git a/xmrstak/backend/cryptonight.hpp b/xmrstak/backend/cryptonight.hpp new file mode 100644 index 0000000..0ef5ae7 --- /dev/null +++ b/xmrstak/backend/cryptonight.hpp @@ -0,0 +1,12 @@ +#pragma once + +// define aeon settings +#define AEON_MEMORY 1048576llu +#define AEON_MASK 0xFFFF0 +#define AEON_ITER 0x40000 + +// define xmr settings +#define MONERO_MEMORY 2097152llu +#define MONERO_MASK 0x1FFFF0 +#define MONERO_ITER 0x80000 + diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp index cf58a21..0bc6214 100644 --- a/xmrstak/backend/nvidia/minethd.cpp +++ b/xmrstak/backend/nvidia/minethd.cpp @@ -32,6 +32,8 @@ #include "xmrstak/jconf.hpp" #include "xmrstak/misc/environment.hpp" #include "xmrstak/backend/cpu/hwlocMemory.hpp" +#include "xmrstak/backend/cryptonight.hpp" +#include "xmrstak/misc/utility.hpp" #include <assert.h> #include <cmath> @@ -208,7 +210,7 @@ void minethd::work_main() uint64_t iCount = 0; cryptonight_ctx* cpu_ctx; cpu_ctx = cpu::minethd::minethd_alloc_ctx(); - cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/); + cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, ::jconf::inst()->IsCurrencyMonero()); uint32_t iNonce; globalStates::inst().iConsumeCnt++; @@ -218,6 +220,9 @@ void minethd::work_main() printer::inst()->print_msg(L0, "Setup failed for GPU %d. Exitting.\n", (int)iThreadNo); std::exit(0); } + + bool mineMonero = strcmp_i(::jconf::inst()->GetCurrency(), "monero"); + bool useAEON = strcmp_i(::jconf::inst()->GetCurrency(), "aeon"); while (bQuit == 0) { @@ -256,7 +261,18 @@ void minethd::work_main() uint32_t foundCount; cryptonight_extra_cpu_prepare(&ctx, iNonce); - cryptonight_core_cpu_hash(&ctx); +#ifndef CONF_NO_MONERO + if(mineMonero) + { + cryptonight_core_cpu_hash<MONERO_ITER, MONERO_MASK, 19>(&ctx); + } +#endif +#ifndef CONF_NO_AEON + if(useAEON) + { + cryptonight_core_cpu_hash<MONERO_ITER, MONERO_MASK, 18>(&ctx); + } +#endif cryptonight_extra_cpu_final(&ctx, iNonce, oWork.iTarget, &foundCount, foundNonce); for(size_t i = 0; i < foundCount; i++) diff --git a/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp b/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp index 784c38d..4e0ace7 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp +++ b/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp @@ -41,8 +41,9 @@ int cuda_get_deviceinfo(nvid_ctx *ctx); int cryptonight_extra_cpu_init(nvid_ctx *ctx); void cryptonight_extra_cpu_set_data( nvid_ctx* ctx, const void *data, uint32_t len); void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce); -void cryptonight_core_cpu_hash(nvid_ctx* ctx); void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce); - } +template<size_t ITERATIONS, size_t THREAD_SHIFT, uint32_t MASK> +void cryptonight_core_cpu_hash(nvid_ctx* ctx); + diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu index e1b78ce..5bea230 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu @@ -92,6 +92,7 @@ __device__ __forceinline__ void storeGlobal32( T* addr, T const & val ) asm volatile( "st.global.cg.u32 [%0], %1;" : : "l"( addr ), "r"( val ) ); } +template<size_t ITERATIONS, size_t THREAD_SHIFT> __global__ void cryptonight_core_gpu_phase1( int threads, int bfactor, int partidx, uint32_t * __restrict__ long_state, uint32_t * __restrict__ ctx_state, uint32_t * __restrict__ ctx_key1 ) { __shared__ uint32_t sharedMemory[1024]; @@ -102,7 +103,7 @@ __global__ void cryptonight_core_gpu_phase1( int threads, int bfactor, int parti const int thread = ( blockDim.x * blockIdx.x + threadIdx.x ) >> 3; const int sub = ( threadIdx.x & 7 ) << 2; - const int batchsize = 0x80000 >> bfactor; + const int batchsize = ITERATIONS >> bfactor; const int start = partidx * batchsize; const int end = start + batchsize; @@ -121,13 +122,13 @@ __global__ void cryptonight_core_gpu_phase1( int threads, int bfactor, int parti else { // load previous text data - MEMCPY8( text, &long_state[( (uint64_t) thread << 19 ) + sub + start - 32], 2 ); + MEMCPY8( text, &long_state[( (uint64_t) thread << THREAD_SHIFT ) + sub + start - 32], 2 ); } __syncthreads( ); for ( int i = start; i < end; i += 32 ) { cn_aes_pseudo_round_mut( sharedMemory, text, key ); - MEMCPY8(&long_state[((uint64_t) thread << 19) + (sub + i)], text, 2); + MEMCPY8(&long_state[((uint64_t) thread << THREAD_SHIFT) + (sub + i)], text, 2); } } @@ -167,6 +168,7 @@ __forceinline__ __device__ uint32_t shuffle(volatile uint32_t* ptr,const uint32_ #ifdef XMR_STAK_THREADS __launch_bounds__( XMR_STAK_THREADS * 4 ) #endif +template<size_t ITERATIONS, size_t THREAD_SHIFT, uint32_t MASK> __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int partidx, uint32_t * d_long_state, uint32_t * d_ctx_a, uint32_t * d_ctx_b ) { __shared__ uint32_t sharedMemory[1024]; @@ -190,10 +192,10 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti int i, k; uint32_t j; - const int batchsize = ITER >> ( 2 + bfactor ); + const int batchsize = (ITERATIONS * 2) >> ( 2 + bfactor ); const int start = partidx * batchsize; const int end = start + batchsize; - uint32_t * long_state = &d_long_state[(IndexType) thread << 19]; + uint32_t * long_state = &d_long_state[(IndexType) thread << THREAD_SHIFT]; uint32_t * ctx_a = d_ctx_a + thread * 4; uint32_t * ctx_b = d_ctx_b + thread * 4; uint32_t a, d[2]; @@ -207,7 +209,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti #pragma unroll 2 for ( int x = 0; x < 2; ++x ) { - j = ( ( shuffle(sPtr,sub, a, 0) & 0x1FFFF0 ) >> 2 ) + sub; + j = ( ( shuffle(sPtr,sub, a, 0) & MASK ) >> 2 ) + sub; const uint32_t x_0 = loadGlobal32<uint32_t>( long_state + j ); const uint32_t x_1 = shuffle(sPtr,sub, x_0, sub + 1); @@ -225,8 +227,8 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti //long_state[j] = d[0] ^ d[1]; storeGlobal32( long_state + j, d[0] ^ d[1] ); - //MUL_SUM_XOR_DST(c, a, &long_state[((uint32_t *)c)[0] & 0x1FFFF0]); - j = ( ( *t1 & 0x1FFFF0 ) >> 2 ) + sub; + //MUL_SUM_XOR_DST(c, a, &long_state[((uint32_t *)c)[0] & MASK]); + j = ( ( *t1 & MASK ) >> 2 ) + sub; uint32_t yy[2]; *( (uint64_t*) yy ) = loadGlobal64<uint64_t>( ( (uint64_t *) long_state )+( j >> 1 ) ); @@ -255,6 +257,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti } } +template<size_t ITERATIONS, size_t THREAD_SHIFT> __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int partidx, const uint32_t * __restrict__ long_state, uint32_t * __restrict__ d_ctx_state, uint32_t * __restrict__ d_ctx_key2 ) { __shared__ uint32_t sharedMemory[1024]; @@ -265,7 +268,7 @@ __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int parti int thread = ( blockDim.x * blockIdx.x + threadIdx.x ) >> 3; int sub = ( threadIdx.x & 7 ) << 2; - const int batchsize = 0x80000 >> bfactor; + const int batchsize = ITERATIONS >> bfactor; const int start = partidx * batchsize; const int end = start + batchsize; @@ -281,7 +284,7 @@ __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int parti { #pragma unroll for ( int j = 0; j < 4; ++j ) - text[j] ^= long_state[((IndexType) thread << 19) + (sub + i + j)]; + text[j] ^= long_state[((IndexType) thread << THREAD_SHIFT) + (sub + i + j)]; cn_aes_pseudo_round_mut( sharedMemory, text, key ); } @@ -289,7 +292,8 @@ __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int parti MEMCPY8( d_ctx_state + thread * 50 + sub + 16, text, 2 ); } -extern "C" void cryptonight_core_cpu_hash(nvid_ctx* ctx) +template<size_t ITERATIONS, size_t THREAD_SHIFT, uint32_t MASK> +void cryptonight_core_cpu_hash(nvid_ctx* ctx) { dim3 grid( ctx->device_blocks ); dim3 block( ctx->device_threads ); @@ -311,7 +315,7 @@ extern "C" void cryptonight_core_cpu_hash(nvid_ctx* ctx) for ( int i = 0; i < partcountOneThree; i++ ) { - CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase1<<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads, + CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase1<ITERATIONS,THREAD_SHIFT><<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads, bfactorOneThree, i, ctx->d_long_state, ctx->d_ctx_state, ctx->d_ctx_key1 )); @@ -321,7 +325,7 @@ extern "C" void cryptonight_core_cpu_hash(nvid_ctx* ctx) for ( int i = 0; i < partcount; i++ ) { - CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase2<<< + CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase2<ITERATIONS,THREAD_SHIFT,MASK><<< grid, block4, block4.x * sizeof(uint32_t) * static_cast< int >( ctx->device_arch[0] < 3 ) @@ -339,7 +343,7 @@ extern "C" void cryptonight_core_cpu_hash(nvid_ctx* ctx) for ( int i = 0; i < partcountOneThree; i++ ) { - CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase3<<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads, + CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase3<ITERATIONS,THREAD_SHIFT><<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads, bfactorOneThree, i, ctx->d_long_state, ctx->d_ctx_state, ctx->d_ctx_key2 )); diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu index 7734473..abca489 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu @@ -5,6 +5,7 @@ #include <cuda_runtime.h> #include <device_functions.hpp> #include <algorithm> +#include "xmrstak/jconf.hpp" #ifdef __CUDACC__ __constant__ @@ -188,8 +189,18 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx) cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); + size_t hashMemSize; + if(::jconf::inst()->IsCurrencyMonero()) + { + hashMemSize = MONERO_MEMORY; + } + else + { + hashMemSize = AEON_MEMORY; + } + size_t wsize = ctx->device_blocks * ctx->device_threads; - CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_long_state, (size_t)MEMORY * wsize)); + CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_long_state, hashMemSize * wsize)); CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_state, 50 * sizeof(uint32_t) * wsize)); CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_key1, 40 * sizeof(uint32_t) * wsize)); CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_key2, 40 * sizeof(uint32_t) * wsize)); @@ -343,13 +354,23 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) ctx->total_device_memory = totalMemory; ctx->free_device_memory = freeMemory; + size_t hashMemSize; + if(::jconf::inst()->IsCurrencyMonero()) + { + hashMemSize = MONERO_MEMORY; + } + else + { + hashMemSize = AEON_MEMORY; + } + // keep 64MiB memory free (value is randomly chosen) // 200byte are meta data memory (result nonce, ...) size_t availableMem = freeMemory - (64u * 1024 * 1024) - 200u; size_t limitedMemory = std::min(availableMem, maxMemUsage); // up to 920bytes extra memory is used per thread for some kernel (lmem/local memory) // 680bytes are extra meta data memory per hash - size_t perThread = size_t(MEMORY) + 740u + 680u; + size_t perThread = hashMemSize + 740u + 680u; size_t max_intensity = limitedMemory / perThread; ctx->device_threads = max_intensity / ctx->device_blocks; // use only odd number of threads diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.hpp b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.hpp index 3ccdcd6..055a8bd 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.hpp +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.hpp @@ -1,5 +1,7 @@ #pragma once +#include "xmrstak/backend/cryptonight.hpp" + #ifdef __INTELLISENSE__ #define __CUDA_ARCH__ 520 /* avoid red underlining */ @@ -18,8 +20,6 @@ struct uint3 blockDim; #define __shfl(a,b,c) 1 #endif -#define MEMORY (1 << 21) // 2 MiB / 2097152 B -#define ITER (1 << 20) // 1048576 #define AES_BLOCK_SIZE 16 #define AES_KEY_SIZE 32 #define INIT_SIZE_BLK 8 diff --git a/xmrstak/cli/cli-miner.cpp b/xmrstak/cli/cli-miner.cpp index 8939803..f571ad6 100644 --- a/xmrstak/cli/cli-miner.cpp +++ b/xmrstak/cli/cli-miner.cpp @@ -31,6 +31,7 @@ #include "xmrstak/params.hpp" #include "xmrstak/misc/configEditor.hpp" #include "xmrstak/version.hpp" +#include "xmrstak/misc/utility.hpp" #ifndef CONF_NO_HTTPD # include "xmrstak/http/httpd.hpp" @@ -62,8 +63,11 @@ void help() cout<<"Usage: "<<params::inst().binaryName<<" [OPTION]..."<<endl; cout<<" "<<endl; - cout<<" -c, --config FILE common miner configuration file"<<endl; cout<<" -h, --help show this help"<<endl; + cout<<" -c, --config FILE common miner configuration file"<<endl; +#if (!defined(CONF_NO_AEON)) && (!defined(CONF_NO_MONERO)) + cout<<" --currency NAME currency to mine: monero or aeon"<<endl; +#endif #ifndef CONF_NO_CPU cout<<" --noCPU disable the CPU miner backend"<<endl; cout<<" --cpu FILE CPU backend miner config file"<<endl; @@ -168,6 +172,17 @@ int main(int argc, char *argv[]) } params::inst().configFileNVIDIA = argv[i]; } + else if(opName.compare("--currency") == 0) + { + ++i; + if( i >=argc ) + { + printer::inst()->print_msg(L0, "No argument for parameter '--currency' given"); + win_exit(); + return 1; + } + params::inst().currency = argv[i]; + } else if(opName.compare("-o") == 0 || opName.compare("--url") == 0) { ++i; @@ -230,10 +245,30 @@ int main(int argc, char *argv[]) ; configEditor configTpl{}; configTpl.set(std::string(tpl)); + std::cout<<"Please enter:"<<std::endl; + auto& currency = params::inst().currency; + if(currency.empty()) + { + std::string tmp; +#if defined(CONF_NO_AEON) + tmp = "monero"; +#elif defined(CONF_NO_MONERO) + tmp = "aeon"; +#endif + while(!xmrstak::strcmp_i(tmp, "monero") && !xmrstak::strcmp_i(tmp, "aeon")) + { + std::cout<<"- currency: 'monero' or 'aeon'"<<std::endl; + std::cin >> tmp; + } + currency = tmp; + } auto& pool = params::inst().poolURL; if(pool.empty()) { - std::cout<<"Please enter:\n- pool address: e.g. pool.usxmrpool.com:3333"<<std::endl; + if(xmrstak::strcmp_i(currency, "monero")) + std::cout<<"- pool address: e.g. pool.usxmrpool.com:3333"<<std::endl; + else + std::cout<<"- pool address: e.g. mine.aeon-pool.com:5555"<<std::endl; std::cin >> pool; } auto& userName = params::inst().poolUsername; @@ -253,6 +288,7 @@ int main(int argc, char *argv[]) configTpl.replace("POOLURL", pool); configTpl.replace("POOLUSER", userName); configTpl.replace("POOLPASSWD", passwd); + configTpl.replace("CURRENCY", currency); configTpl.write(params::inst().configFile); std::cout<<"Configuration stored in file '"<<params::inst().configFile<<"'"<<std::endl; } @@ -298,6 +334,10 @@ int main(int argc, char *argv[]) printer::inst()->print_str("'r' - results\n"); printer::inst()->print_str("'c' - connection\n"); printer::inst()->print_str("-------------------------------------------------------------------\n"); + if(::jconf::inst()->IsCurrencyMonero()) + printer::inst()->print_msg(L0,"Start mining: MONERO"); + else + printer::inst()->print_msg(L0,"Start mining: AEON"); if(strlen(jconf::inst()->GetOutputFile()) != 0) printer::inst()->open_logfile(jconf::inst()->GetOutputFile()); diff --git a/xmrstak/config.tpl b/xmrstak/config.tpl index 60d85cd..032d483 100644 --- a/xmrstak/config.tpl +++ b/xmrstak/config.tpl @@ -11,6 +11,12 @@ R"===( "pool_password" : "POOLPASSWD", /* + * currency to mine + * allowed values: 'monero' or 'aeon' + */ +"currency" : "CURRENCY", + +/* * Network timeouts. * Because of the way this client is written it doesn't need to constantly talk (keep-alive) to the server to make * sure it is there. We detect a buggy / overloaded server by the call timeout. The default values will be ok for diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp index c033e66..4b23ed5 100644 --- a/xmrstak/jconf.cpp +++ b/xmrstak/jconf.cpp @@ -27,6 +27,7 @@ #include "xmrstak/misc/console.hpp" #include "xmrstak/misc/jext.hpp" #include "xmrstak/misc/console.hpp" +#include "xmrstak/misc/utility.hpp" #include <stdio.h> #include <stdlib.h> @@ -46,7 +47,7 @@ using namespace rapidjson; * This enum needs to match index in oConfigValues, otherwise we will get a runtime error */ enum configEnum { - bTlsMode, bTlsSecureAlgo, sTlsFingerprint, sPoolAddr, sWalletAddr, sPoolPwd, + bTlsMode, bTlsSecureAlgo, sTlsFingerprint, sPoolAddr, sWalletAddr, sPoolPwd,sCurrency, iCallTimeout, iNetRetry, iGiveUpLimit, iVerboseLevel, iAutohashTime, bDaemonMode, sOutputFile, iHttpdPort, bPreferIpv4, bNiceHashMode, bAesOverride, sUseSlowMem }; @@ -65,6 +66,7 @@ configVal oConfigValues[] = { { sPoolAddr, "pool_address", kStringType }, { sWalletAddr, "wallet_address", kStringType }, { sPoolPwd, "pool_password", kStringType }, + { sCurrency, "currency", kStringType }, { iCallTimeout, "call_timeout", kNumberType }, { iNetRetry, "retry_time", kNumberType }, { iGiveUpLimit, "giveup_limit", kNumberType }, @@ -150,6 +152,45 @@ const char* jconf::GetWalletAddress() return poolUsername.c_str(); } +const std::string jconf::GetCurrency() +{ + auto& currency = xmrstak::params::inst().currency; + if(currency.empty()) + currency = prv->configValues[sCurrency]->GetString(); + if( +#ifndef CONF_NO_MONERO + // if monero is disabled at compile time, enable error message if selected currency is `monero` + !xmrstak::strcmp_i(currency, "monero") +#else + true +#endif + && +#ifndef CONF_NO_AEON + // if aeon is disabled at compile time, enable error message if selected currency is `aeon` + !xmrstak::strcmp_i(currency, "aeon") +#else + true +#endif + ) + { + printer::inst()->print_msg(L0, "ERROR: Wrong currency selected - '%s'.", currency.c_str()); + win_exit(); + } + return currency; +} + +bool jconf::IsCurrencyMonero() +{ + if(xmrstak::strcmp_i(::jconf::inst()->GetCurrency(), "monero")) + { + return true; + } + else + { + return false; + } +} + bool jconf::PreferIpv4() { return prv->configValues[bPreferIpv4]->GetBool(); diff --git a/xmrstak/jconf.hpp b/xmrstak/jconf.hpp index 688ffe1..48b47b5 100644 --- a/xmrstak/jconf.hpp +++ b/xmrstak/jconf.hpp @@ -41,6 +41,8 @@ public: const char* GetPoolAddress(); const char* GetPoolPwd(); const char* GetWalletAddress(); + const std::string GetCurrency(); + bool IsCurrencyMonero(); uint64_t GetVerboseLevel(); uint64_t GetAutohashTime(); diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp index 7fc46e4..b469dc2 100644 --- a/xmrstak/misc/executor.cpp +++ b/xmrstak/misc/executor.cpp @@ -21,6 +21,7 @@ * */ +#include "xmrstak/jconf.hpp" #include "executor.hpp" #include "xmrstak/net/jpsock.hpp" @@ -183,8 +184,16 @@ void executor::on_sock_ready(size_t pool_id) if(pool_id == dev_pool_id) { - if(!pool->cmd_login("", "")) - pool->disconnect(); + if(::jconf::inst()->IsCurrencyMonero()) + { + if(!pool->cmd_login("", "")) + pool->disconnect(); + } + else + { + if(!pool->cmd_login("WmsvqXDu7Fw5eAEZr1euJH3ycad55NxFd82PfhLR9Zi1Nq5S74zk63EA8fyMS8BQNR94os9N9aah87inKkumNJ7G2d7qTpRLN", "x")) + pool->disconnect(); + } current_pool_id = dev_pool_id; printer::inst()->print_msg(L1, "Dev pool logged in. Switching work."); @@ -352,8 +361,12 @@ void executor::on_switch_pool(size_t pool_id) // If it fails, it fails, we carry on on the usr pool // as we never receive further events printer::inst()->print_msg(L1, "Connecting to dev pool..."); - const char* dev_pool_addr = jconf::inst()->GetTlsSetting() ? "donate.xmr-stak.net:6666" : "donate.xmr-stak.net:3333"; - if(!pool->connect(dev_pool_addr, error)) + std::string dev_pool_addr; + if(::jconf::inst()->IsCurrencyMonero()) + dev_pool_addr = jconf::inst()->GetTlsSetting() ? "donate.xmr-stak.net:6666" : "donate.xmr-stak.net:3333"; + else + dev_pool_addr = jconf::inst()->GetTlsSetting() ? "mine.aeon-pool.com:443" : "mine.aeon-pool.com:5555"; + if(!pool->connect(dev_pool_addr.c_str(), error)) printer::inst()->print_msg(L1, "Error connecting to dev pool. Staying with user pool."); } else diff --git a/xmrstak/misc/utility.cpp b/xmrstak/misc/utility.cpp new file mode 100644 index 0000000..3b1369a --- /dev/null +++ b/xmrstak/misc/utility.cpp @@ -0,0 +1,21 @@ +#include <string> +#include <algorithm> + + +namespace xmrstak +{ + bool strcmp_i(const std::string& str1, const std::string& str2) + { + if(str1.size() != str2.size()) + return false; + else + return (str1.empty() | str2.empty()) ? + false : + std::equal(str1.begin(), str1.end(),str2.begin(), + [](char c1, char c2) + { + return ::tolower(c1) == ::tolower(c2); + } + ); + } +} // namepsace xmrstak diff --git a/xmrstak/misc/utility.hpp b/xmrstak/misc/utility.hpp new file mode 100644 index 0000000..b2e841d --- /dev/null +++ b/xmrstak/misc/utility.hpp @@ -0,0 +1,12 @@ +#pragma once + +#include <string> + +namespace xmrstak +{ + /** case insensitive string compare + * + * @return true if both strings are equal, else false + */ + bool strcmp_i(const std::string& str1, const std::string& str2); +} // namepsace xmrstak diff --git a/xmrstak/params.hpp b/xmrstak/params.hpp index dddf82e..6127212 100644 --- a/xmrstak/params.hpp +++ b/xmrstak/params.hpp @@ -28,6 +28,8 @@ struct params std::string poolPasswd; std::string poolUsername; + std::string currency; + std::string configFile; std::string configFileAMD; std::string configFileNVIDIA; |
