diff options
37 files changed, 533 insertions, 243 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt index 10f33bd..3b3c7eb 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,6 +1,6 @@ project(xmr-stak) -cmake_minimum_required(VERSION 3.1.0) +cmake_minimum_required(VERSION 3.4.0) # enforce C++11 set(CMAKE_CXX_STANDARD_REQUIRED ON) @@ -522,9 +522,11 @@ endif() file(GLOB SRCFILES_CPP "xmrstak/cli/*.cpp") set_source_files_properties(${SRCFILES_CPP} PROPERTIES LANGUAGE CXX) -add_executable(xmr-stak - ${SRCFILES_CPP} -) +if(CMAKE_CXX_COMPILER_ID MATCHES "MSVC") + add_executable(xmr-stak ${SRCFILES_CPP} xmrstak/cli/xmr-stak.manifest) +else() + add_executable(xmr-stak ${SRCFILES_CPP}) +endif() set(EXECUTABLE_OUTPUT_PATH "bin") set(LIBRARY_OUTPUT_PATH "bin") @@ -8,7 +8,7 @@ ENV XMRSTAK_CMAKE_FLAGS -DXMR-STAK_COMPILE=generic -DCUDA_ENABLE=ON -DOpenCL_ENA # Innstall packages RUN apt-get update \ && set -x \ - && apt-get install -qq --no-install-recommends -y ca-certificates cmake cuda-core-9-0 git cuda-cudart-dev-9-0 libhwloc-dev libmicrohttpd-dev libssl-dev \ + && apt-get install -qq --no-install-recommends -y build-essential ca-certificates cmake cuda-core-9-0 git cuda-cudart-dev-9-0 libhwloc-dev libmicrohttpd-dev libssl-dev \ && git clone $GIT_REPOSITORY \ && cd /xmr-stak \ && cmake ${XMRSTAK_CMAKE_FLAGS} . \ @@ -16,7 +16,7 @@ RUN apt-get update \ && cd - \ && mv /xmr-stak/bin/* /usr/local/bin/ \ && rm -rf /xmr-stak \ - && apt-get purge -y -qq cmake cuda-core-9-0 git cuda-cudart-dev-9-0 libhwloc-dev libmicrohttpd-dev libssl-dev \ + && apt-get purge -y -qq build-essential cmake cuda-core-9-0 git cuda-cudart-dev-9-0 libhwloc-dev libmicrohttpd-dev libssl-dev \ && apt-get clean -qq VOLUME /mnt @@ -21,15 +21,15 @@ XMR-Stak is a universal Stratum pool miner. This miner supports CPUs, AMD and NV ## Features - support all common backends (CPU/x86, AMD-GPU and NVIDIA-GPU) -- support all common OS (Linux, Windows and MacOS) +- 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 - open source software (GPLv3) - TLS support -- HTML statistics -- JSON API for monitoring +- [HTML statistics](doc/usage.md#html-and-json-api-report-configuraton) +- [JSON API for monitoring](doc/usage.md#html-and-json-api-report-configuraton) ## Supported altcoins @@ -62,7 +62,7 @@ fireice-uk: psychocrypt: ``` -43NoJVEXo21hGZ6tDG6Z3g4qimiGdJPE6GRxAmiWwm26gwr62Lqo7zRiCJFSBmbkwTGNuuES9ES5TgaVHceuYc4Y75txCTU +45tcqnJMgd3VqeTznNotiNj4G9PQoK67TGRiHyj6EYSZ31NUbAfs9XdiU5squmZb717iHJLxZv3KfEw8jCYGL5wa19yrVCn ``` ## Release Checksums @@ -1,7 +1,7 @@ # FAQ ## Content Overview -* [SeLockMemoryPrivilege failed](#selockmemoryprivilege-failed) +* ["Obtaining SeLockMemoryPrivilege failed."](#obtaining-selockmemoryprivilege-failed) * [VirtualAlloc failed](#virtualalloc-failed) * [Error msvcp140.dll and vcruntime140.dll not available](#error-msvcp140dll-and-vcruntime140dll-not-available) * [Error: MEMORY ALLOC FAILED: mmap failed](#error-memory-alloc-failed-mmap-failed) @@ -9,23 +9,24 @@ * [Virus Protection Alert](#virus-protection-alert) * [Change Currency to Mine](#change-currency-to-mine) -## SeLockMemoryPrivilege failed +## "Obtaining SeLockMemoryPrivilege failed." -Please see [config.txt](config.txt) under section **LARGE PAGE SUPPORT** +For professional versions of Windows see [this article](https://msdn.microsoft.com/en-gb/library/ms190730.aspx). +Make sure to reboot afterwards! -For Windows 7 pro, or Windows 8 and above see [this article](https://msdn.microsoft.com/en-gb/library/ms190730.aspx) (make sure to reboot afterwards!). +For Windows 7/10 Home: -For Windows 7 Home : +1) Download and install [Windows Server 2003 Resource Kit Tools](https://www.microsoft.com/en-us/download/details.aspx?id=17657). Ignore any incompatibility warning during installation. -1) Download and install [Windows Server 2003 Resource Kit Tools](https://www.microsoft.com/en-us/download/details.aspx?id=17657). Ignore incompatiablity warning during installation. +2) Open cmd or PowerShell as an administrator. -2) In cmd or power shell: `ntrights -u %USERNAME% +r SeLockMemoryPrivilege` (where %USERNAME% is the user that will be running the program. This command needs to be run as admin) +3) Use `ntrights -u %USERNAME% +r SeLockMemoryPrivilege` where %USERNAME% is the user that will be running the program. -3) Reboot. +4) Reboot. Reference: http://rybkaforum.net/cgi-bin/rybkaforum/topic_show.pl?pid=259791#pid259791 -*Warning: do not download ntrights.exe from any other site other then the offical Microsoft download page.* +*Warning: Do not download ntrights.exe from any other site other than the offical Microsoft download page.* ## VirtualAlloc failed @@ -40,13 +41,18 @@ Download and install this [runtime package](https://go.microsoft.com/fwlink/?Lin ## Error: MEMORY ALLOC FAILED: mmap failed -On Linux you will need to configure large page support `sudo sysctl -w vm.nr_hugepages=128` and increase your -ulimit -l. To do this you need to add following lines to /etc/security/limits.conf: +On Linux you will need to configure large page support and increase your ulimit -l. + +To set large page support, add the following lines to /etc/sysctl.conf: + + vm.nr_hugepages=128 + +To increase the ulimit, add following lines to /etc/security/limits.conf: * soft memlock 262144 * hard memlock 262144 -Save file. You WILL need to log out and log back in for these settings to take affect on your user (no need to reboot, just relogin in your session). +You WILL need to log out and log back in for these settings to take affect on your user (no need to reboot, just relogin in your session). You can also do it Windows-style and simply run-as-root, but this is NOT recommended for security reasons. diff --git a/doc/Linux_deployment.md b/doc/Linux_deployment.md index 323a97f..3219e8a 100644 --- a/doc/Linux_deployment.md +++ b/doc/Linux_deployment.md @@ -15,6 +15,7 @@ For automatic deployments, please use the steps above to obtain config.txt and u ``` #!/bin/bash +sudo apt install curl curl -O `curl -s https://api.github.com/repos/fireice-uk/xmr-stak/releases/latest | grep -o 'browser_download_url.*xmr-stak-portbin-linux.tar.gz' | sed 's/.*"//'` curl -O http://path.to/your/config.txt tar xzf xmr-stak-portbin-linux.tar.gz diff --git a/doc/compile.md b/doc/compile.md index e97affa..771c9d1 100644 --- a/doc/compile.md +++ b/doc/compile.md @@ -9,7 +9,7 @@ * [Compile on Windows](compile_Windows.md) * [Compile on Linux](compile_Linux.md) * [Compile on FreeBSD](compile_FreeBSD.md) -* [Compile on MacOS](compile_MacOS.md) +* [Compile on macOS](compile_macOS.md) ## Build System @@ -31,7 +31,7 @@ After the configuration you need to compile the miner, follow the guide for your * [Compile in Windows](compile_Windows.md) * [Compile in Linux](compile_Linux.md) * [Compile in FreeBSD](compile_FreeBSD.md) -* [Compile in MacOS](compile_MacOS.md) +* [Compile in macOS](compile_macOS.md) ## Generic Build Options - `CMAKE_INSTALL_PREFIX` install miner to the home folder diff --git a/doc/compile_Linux.md b/doc/compile_Linux.md index a3031be..729f4d2 100644 --- a/doc/compile_Linux.md +++ b/doc/compile_Linux.md @@ -25,7 +25,7 @@ make install # Arch - sudo pacman -S base-devel hwloc openssl cmake libmicrohttpd + sudo pacman -S --needed base-devel hwloc openssl cmake libmicrohttpd git clone https://github.com/fireice-uk/xmr-stak.git mkdir xmr-stak/build cd xmr-stak/build @@ -64,6 +64,34 @@ cd xmr-stak/build cmake .. make install + + # TinyCore Linux 8.x + # TinyCore is 32-bit only, but there is an x86-64 port, known as "Pure 64," + # hosted on the TinyCore home page, and it works well. + # Beware that huge page support is not enabled in the kernel distributed + # with Pure 64. Consider http://wiki.tinycorelinux.net/wiki:custom_kernel + # Note that as of yet there are no distro packages for microhttpd or hwloc. + # hwloc is easy enough to install manually though, shown below. + # Also note that only CPU mining has been tested on this platform, thus the + # disabling of CUDA and OpenCL shown below. + tce-load -iw openssl-dev.tcz cmake.tcz make.tcz gcc.tcz git.tcz \ + glibc_base-dev.tcz linux-4.8.1_api_headers.tcz \ + glibc_add_lib.tcz + wget https://www.open-mpi.org/software/hwloc/v1.11/downloads/hwloc-1.11.8.tar.gz + tar xzvf hwloc-1.11.8.tar.gz + cd hwloc-1.11.8 + ./configure --prefix=/usr/local + make + sudo make install + cd .. + git clone http://github.com/fireice-uk/xmr-stak + cd xmr-stak + mkdir build + cd build + CC=gcc cmake .. -DCUDA_ENABLE=OFF \ + -DOpenCL_ENABLE=OFF \ + -DMICROHTTPD_ENABLE=OFF + make install ``` - g++ version 5.1 or higher is required for full C++11 support. diff --git a/doc/compile_Windows.md b/doc/compile_Windows.md index c9a8ff7..d1be866 100644 --- a/doc/compile_Windows.md +++ b/doc/compile_Windows.md @@ -13,6 +13,7 @@ - during the install chose the components - `Desktop development with C++` (left side) - `VC++ 2015.3 v140 toolset for desktop` (right side) + - Since release of VS2017 15.5 (12/04/17), require `VC++ 2017 version 15.4 v14.11 toolset` (under tab `Individual Components`, section `Compilers, build tools, and runtimes`), as CUDA 9.1 is not compatible with compiler 14.12.X ### CMake for Win64 @@ -22,9 +23,9 @@ ### Cuda 8.0+ (only needed to use NVIDIA GPUs) -- donwload and install [https://developer.nvidia.com/cuda-downloads](https://developer.nvidia.com/cuda-downloads) +- download and install [https://developer.nvidia.com/cuda-downloads](https://developer.nvidia.com/cuda-downloads) - for minimal install choose `Custom installation options` during the install and select - - CUDA/Develpment + - CUDA/Development - CUDA/Visual Studio Integration (ignore the warning during the install that VS2017 is not supported) - CUDA/Runtime - Driver components @@ -75,16 +76,21 @@ ## Compile -- download and unzip `xmr-stak` +- download `xmr-stak` [Source Code.zip](https://github.com/fireice-uk/xmr-stak/releases) and save to a location in your home folder (C:\Users\USERNAME\) +- extract the Source Code.zip file - open the command line terminal `cmd` -- `cd` to your unzipped source code directory -- execute the following commands (NOTE: path to VS2017 can be different) +- `cd` to the extracted source code directory +- execute the following commands (NOTE: path to Visual Studio 2017 Community can be different) ``` + # Next line is only if compiling for Cuda 9.1 and using Visual Studio 2017 >= 15.5 (released 12/04/17) + "C:\Program Files (x86)\Microsoft Visual Studio\2017\Community\VC\Auxiliary\Build\vcvarsall.bat" x64 -vcvars_ver=14.11 "C:\Program Files (x86)\Microsoft Visual Studio\2017\Community\Common7\Tools\VsMSBuildCmd.bat" set CMAKE_PREFIX_PATH=C:\xmr-stak-dep\hwloc;C:\xmr-stak-dep\libmicrohttpd;C:\xmr-stak-dep\openssl mkdir build cd build ``` + - Sometimes Windows will `cd` to `C:\Users\USERNAME\source\build\` instead of `C:\Users\USERNAME\xmr-stak-<version>\build`. Ensure you are in the correct `build` directory before proceeding. + - for CUDA 8* ``` cmake -G "Visual Studio 15 2017 Win64" -T v140,host=x64 .. diff --git a/doc/compile_MacOS.md b/doc/compile_macOS.md index 1b0af91..46f1d5b 100644 --- a/doc/compile_MacOS.md +++ b/doc/compile_macOS.md @@ -1,4 +1,4 @@ -# Compile **xmr-stak** for MacOS +# Compile **xmr-stak** for macOS ## Dependencies @@ -18,7 +18,13 @@ make install ### For AMD GPUs -> 🖐 We need help with AMD GPU compilation instructions. Please submit a PR if you managed to install [AMD APP SDK](http://developer.amd.com/amd-accelerated-parallel-processing-app-sdk/) and to compile `xmr-stak` on MacOS. +OpenCL is bundled with Xcode, so no other depedency then the basic ones needed. Just enable OpenCL via the `-DOpenCL_ENABLE=ON` CMake option. + +```shell +brew install hwloc libmicrohttpd gcc openssl cmake +cmake . -DOPENSSL_ROOT_DIR=/usr/local/opt/openssl -DCUDA_ENABLE=OFF -DOpenCL_ENABLE=ON +make install +``` ### For CPU-only mining diff --git a/doc/usage.md b/doc/usage.md index 60cf69b..a810469 100644 --- a/doc/usage.md +++ b/doc/usage.md @@ -23,7 +23,7 @@ The number of files depends on the available backends. `set XMRSTAK_NOWAIT=1` disable the dialog `Press any key to exit.` for non UAC execution. -## Usage on Linux & MacOS +## Usage on Linux & macOS 1) Open a terminal within the folder with the binary 2) Start the miner with `./xmr-stak` diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index d9bc962..c39c567 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -518,13 +518,13 @@ std::vector<GpuContext> getAMDDevices(int index) printer::inst()->print_msg(L1,"WARNING: %s when calling clGetDeviceInfo to get CL_DEVICE_NAME for device %u.", err_to_str(clStatus), k); continue; } - printer::inst()->print_msg(L0,"Found OpenCL GPU %s.",ctx.name.c_str()); // if environment variable GPU_SINGLE_ALLOC_PERCENT is not set we can not allocate the full memory ctx.deviceIdx = k; ctx.freeMem = std::min(ctx.freeMem, maxMem); ctx.name = std::string(devNameVec.data()); ctx.DeviceID = device_list[k]; + printer::inst()->print_msg(L0,"Found OpenCL GPU %s.",ctx.name.c_str()); ctxVec.push_back(ctx); } } diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl index 255fcbb..ec05712 100644 --- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl +++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl @@ -653,21 +653,11 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u for(int i = 0; i < 25; ++i) states[i] = State[i]; - switch(State[0] & 3) - { - case 0: - Branch0[atomic_inc(Branch0 + Threads)] = get_global_id(0) - get_global_offset(0); - break; - case 1: - Branch1[atomic_inc(Branch1 + Threads)] = get_global_id(0) - get_global_offset(0); - break; - case 2: - Branch2[atomic_inc(Branch2 + Threads)] = get_global_id(0) - get_global_offset(0); - break; - case 3: - Branch3[atomic_inc(Branch3 + Threads)] = get_global_id(0) - get_global_offset(0); - break; - } + ulong StateSwitch = State[0] & 3; + __global uint *destinationBranch1 = StateSwitch == 0 ? Branch0 : Branch1; + __global uint *destinationBranch2 = StateSwitch == 2 ? Branch2 : Branch3; + __global uint *destinationBranch = StateSwitch < 2 ? destinationBranch1 : destinationBranch2; + destinationBranch[atomic_inc(destinationBranch + Threads)] = gIdx; } } mem_fence(CLK_GLOBAL_MEM_FENCE); @@ -704,8 +694,7 @@ __kernel void Skein(__global ulong *states, __global uint *BranchBuf, __global u for(uint i = 0; i < 4; ++i) { - if(i < 3) t[0] += 0x40UL; - else t[0] += 0x08UL; + t[0] += i < 3 ? 0x40UL : 0x08UL; t[2] = t[0] ^ t[1]; @@ -715,8 +704,7 @@ __kernel void Skein(__global ulong *states, __global uint *BranchBuf, __global u h = m ^ p; - if(i < 2) t[1] = 0x3000000000000000UL; - else t[1] = 0xB000000000000000UL; + t[1] = i < 2 ? 0x3000000000000000UL : 0xB000000000000000UL; } t[0] = 0x08UL; @@ -744,6 +732,27 @@ __kernel void Skein(__global ulong *states, __global uint *BranchBuf, __global u #define SWAP8(x) as_ulong(as_uchar8(x).s76543210) +#define JHXOR \ + h0h ^= input[0]; \ + h0l ^= input[1]; \ + h1h ^= input[2]; \ + h1l ^= input[3]; \ + h2h ^= input[4]; \ + h2l ^= input[5]; \ + h3h ^= input[6]; \ + h3l ^= input[7]; \ +\ + E8; \ +\ + h4h ^= input[0]; \ + h4l ^= input[1]; \ + h5h ^= input[2]; \ + h5l ^= input[3]; \ + h6h ^= input[4]; \ + h6l ^= input[5]; \ + h7h ^= input[6]; \ + h7l ^= input[7] + __kernel void JH(__global ulong *states, __global uint *BranchBuf, __global uint *output, ulong Target, ulong Threads) { const uint idx = get_global_id(0) - get_global_offset(0); @@ -757,46 +766,27 @@ __kernel void JH(__global ulong *states, __global uint *BranchBuf, __global uint sph_u64 h4h = 0x754D2E7F8996A371UL, h4l = 0x62E27DF70849141DUL, h5h = 0x948F2476F7957627UL, h5l = 0x6C29804757B6D587UL, h6h = 0x6C0D8EAC2D275E5CUL, h6l = 0x0F7A0557C6508451UL, h7h = 0xEA12247067D3E47BUL, h7l = 0x69D71CD313ABE389UL; sph_u64 tmp; - for(int i = 0; i < 5; ++i) + for(int i = 0; i < 3; ++i) { ulong input[8]; - if(i < 3) - { - for(int x = 0; x < 8; ++x) input[x] = (states[(i << 3) + x]); - } - else if(i == 3) - { - input[0] = (states[24]); - input[1] = 0x80UL; - for(int x = 2; x < 8; ++x) input[x] = 0x00UL; - } - else - { - input[7] = 0x4006000000000000UL; - - for(int x = 0; x < 7; ++x) input[x] = 0x00UL; - } - - h0h ^= input[0]; - h0l ^= input[1]; - h1h ^= input[2]; - h1l ^= input[3]; - h2h ^= input[4]; - h2l ^= input[5]; - h3h ^= input[6]; - h3l ^= input[7]; - - E8; - - h4h ^= input[0]; - h4l ^= input[1]; - h5h ^= input[2]; - h5l ^= input[3]; - h6h ^= input[4]; - h6l ^= input[5]; - h7h ^= input[6]; - h7l ^= input[7]; + const int shifted = i << 3; + for(int x = 0; x < 8; ++x) input[x] = (states[shifted + x]); + JHXOR; + } + { + ulong input[8]; + input[0] = (states[24]); + input[1] = 0x80UL; + #pragma unroll 6 + for(int x = 2; x < 8; ++x) input[x] = 0x00UL; + JHXOR; + } + { + ulong input[8]; + for(int x = 0; x < 7; ++x) input[x] = 0x00UL; + input[7] = 0x4006000000000000UL; + JHXOR; } //output[0] = h6h; @@ -832,6 +822,7 @@ __kernel void Blake(__global ulong *states, __global uint *BranchBuf, __global u ((uint8 *)h)[0] = vload8(0U, c_IV256); + #pragma unroll 4 for(uint i = 0, bitlen = 0; i < 4; ++i) { if(i < 3) @@ -907,6 +898,7 @@ __kernel void Groestl(__global ulong *states, __global uint *BranchBuf, __global State[7] = 0x0001000000000000UL; + #pragma unroll 4 for(uint i = 0; i < 4; ++i) { ulong H[8], M[8]; diff --git a/xmrstak/backend/amd/amd_gpu/opencl/wolf-skein.cl b/xmrstak/backend/amd/amd_gpu/opencl/wolf-skein.cl index 868757b..279b652 100644 --- a/xmrstak/backend/amd/amd_gpu/opencl/wolf-skein.cl +++ b/xmrstak/backend/amd/amd_gpu/opencl/wolf-skein.cl @@ -3,6 +3,7 @@ R"===( #define WOLF_SKEIN_CL // Vectorized Skein implementation macros and functions by Wolf +// Updated by taisel #define SKEIN_KS_PARITY 0x1BD11BDAA9FC1A22 @@ -22,11 +23,11 @@ static const __constant ulong SKEIN512_256_IV[8] = 0xC36FBAF9393AD185UL, 0x3EEDBA1833EDFC13UL }; -#define SKEIN_INJECT_KEY(p, s) do { \ +#define SKEIN_INJECT_KEY(p, s, q) do { \ p += h; \ - p.s5 += t[s % 3]; \ - p.s6 += t[(s + 1) % 3]; \ - p.s7 += s; \ + p.s5 += t[s]; \ + p.s6 += t[select(s + 1U, 0U, s == 2U)]; \ + p.s7 += q; \ } while(0) ulong SKEIN_ROT(const uint2 x, const uint y) @@ -35,55 +36,55 @@ ulong SKEIN_ROT(const uint2 x, const uint y) else return(as_ulong(amd_bitalign(x.s10, x, 32 - (y - 32)))); } -void SkeinMix8(ulong4 *pv0, ulong4 *pv1, const uint rc0, const uint rc1, const uint rc2, const uint rc3) +void SkeinMix8(ulong4 *pv0, ulong4 *pv1, const ulong4 rc) { *pv0 += *pv1; - (*pv1).s0 = SKEIN_ROT(as_uint2((*pv1).s0), rc0); - (*pv1).s1 = SKEIN_ROT(as_uint2((*pv1).s1), rc1); - (*pv1).s2 = SKEIN_ROT(as_uint2((*pv1).s2), rc2); - (*pv1).s3 = SKEIN_ROT(as_uint2((*pv1).s3), rc3); + (*pv1).s0 = SKEIN_ROT(as_uint2((*pv1).s0), rc.s0); + (*pv1).s1 = SKEIN_ROT(as_uint2((*pv1).s1), rc.s1); + (*pv1).s2 = SKEIN_ROT(as_uint2((*pv1).s2), rc.s2); + (*pv1).s3 = SKEIN_ROT(as_uint2((*pv1).s3), rc.s3); *pv1 ^= *pv0; } -ulong8 SkeinEvenRound(ulong8 p, const ulong8 h, const ulong *t, const uint s) +ulong8 SkeinEvenRound(ulong8 p, const ulong8 h, const ulong *t, const uint s, const uint q) { - SKEIN_INJECT_KEY(p, s); + SKEIN_INJECT_KEY(p, s, q); ulong4 pv0 = p.even, pv1 = p.odd; - SkeinMix8(&pv0, &pv1, 46, 36, 19, 37); + SkeinMix8(&pv0, &pv1, (ulong4)(46, 36, 19, 37)); pv0 = shuffle(pv0, (ulong4)(1, 2, 3, 0)); pv1 = shuffle(pv1, (ulong4)(0, 3, 2, 1)); - SkeinMix8(&pv0, &pv1, 33, 27, 14, 42); + SkeinMix8(&pv0, &pv1, (ulong4)(33, 27, 14, 42)); pv0 = shuffle(pv0, (ulong4)(1, 2, 3, 0)); pv1 = shuffle(pv1, (ulong4)(0, 3, 2, 1)); - SkeinMix8(&pv0, &pv1, 17, 49, 36, 39); + SkeinMix8(&pv0, &pv1, (ulong4)(17, 49, 36, 39)); pv0 = shuffle(pv0, (ulong4)(1, 2, 3, 0)); pv1 = shuffle(pv1, (ulong4)(0, 3, 2, 1)); - SkeinMix8(&pv0, &pv1, 44, 9, 54, 56); + SkeinMix8(&pv0, &pv1, (ulong4)(44, 9, 54, 56)); return(shuffle2(pv0, pv1, (ulong8)(1, 4, 2, 7, 3, 6, 0, 5))); } -ulong8 SkeinOddRound(ulong8 p, const ulong8 h, const ulong *t, const uint s) +ulong8 SkeinOddRound(ulong8 p, const ulong8 h, const ulong *t, const uint s, const uint q) { - SKEIN_INJECT_KEY(p, s); + SKEIN_INJECT_KEY(p, s, q); ulong4 pv0 = p.even, pv1 = p.odd; - SkeinMix8(&pv0, &pv1, 39, 30, 34, 24); + SkeinMix8(&pv0, &pv1, (ulong4)(39, 30, 34, 24)); pv0 = shuffle(pv0, (ulong4)(1, 2, 3, 0)); pv1 = shuffle(pv1, (ulong4)(0, 3, 2, 1)); - SkeinMix8(&pv0, &pv1, 13, 50, 10, 17); + SkeinMix8(&pv0, &pv1, (ulong4)(13, 50, 10, 17)); pv0 = shuffle(pv0, (ulong4)(1, 2, 3, 0)); pv1 = shuffle(pv1, (ulong4)(0, 3, 2, 1)); - SkeinMix8(&pv0, &pv1, 25, 29, 39, 43); + SkeinMix8(&pv0, &pv1, (ulong4)(25, 29, 39, 43)); pv0 = shuffle(pv0, (ulong4)(1, 2, 3, 0)); pv1 = shuffle(pv1, (ulong4)(0, 3, 2, 1)); - SkeinMix8(&pv0, &pv1, 8, 35, 56, 22); + SkeinMix8(&pv0, &pv1, (ulong4)(8, 35, 56, 22)); return(shuffle2(pv0, pv1, (ulong8)(1, 4, 2, 7, 3, 6, 0, 5))); } @@ -92,20 +93,47 @@ ulong8 Skein512Block(ulong8 p, ulong8 h, ulong h8, const ulong *t) #pragma unroll for(int i = 0; i < 18; ++i) { - p = SkeinEvenRound(p, h, t, i); + p = SkeinEvenRound(p, h, t, 0U, i); ++i; ulong tmp = h.s0; h = shuffle(h, (ulong8)(1, 2, 3, 4, 5, 6, 7, 0)); h.s7 = h8; h8 = tmp; - p = SkeinOddRound(p, h, t, i); + p = SkeinOddRound(p, h, t, 1U, i); + ++i; + tmp = h.s0; + h = shuffle(h, (ulong8)(1, 2, 3, 4, 5, 6, 7, 0)); + h.s7 = h8; + h8 = tmp; + p = SkeinEvenRound(p, h, t, 2U, i); + ++i; + tmp = h.s0; + h = shuffle(h, (ulong8)(1, 2, 3, 4, 5, 6, 7, 0)); + h.s7 = h8; + h8 = tmp; + p = SkeinOddRound(p, h, t, 0U, i); + ++i; + tmp = h.s0; + h = shuffle(h, (ulong8)(1, 2, 3, 4, 5, 6, 7, 0)); + h.s7 = h8; + h8 = tmp; + p = SkeinEvenRound(p, h, t, 1U, i); + ++i; + tmp = h.s0; + h = shuffle(h, (ulong8)(1, 2, 3, 4, 5, 6, 7, 0)); + h.s7 = h8; + h8 = tmp; + p = SkeinOddRound(p, h, t, 2U, i); tmp = h.s0; h = shuffle(h, (ulong8)(1, 2, 3, 4, 5, 6, 7, 0)); h.s7 = h8; h8 = tmp; } - SKEIN_INJECT_KEY(p, 18); + p += h; + p.s5 += t[0]; + p.s6 += t[1]; + p.s7 += 18; return(p); } diff --git a/xmrstak/backend/amd/autoAdjust.hpp b/xmrstak/backend/amd/autoAdjust.hpp index 0bc5239..afedb5c 100644 --- a/xmrstak/backend/amd/autoAdjust.hpp +++ b/xmrstak/backend/amd/autoAdjust.hpp @@ -94,14 +94,22 @@ private: } std::string conf; - int i = 0; for(auto& ctx : devVec) { /* 1000 is a magic selected limit, the reason is that more than 2GiB memory * sowing down the memory performance because of TLB cache misses */ size_t maxThreads = 1000u; - if(ctx.name.compare("gfx901") == 0) + if( + ctx.name.compare("gfx901") == 0 || + ctx.name.compare("gfx904") == 0 || + // APU + ctx.name.compare("gfx902") == 0 || + // UNKNOWN + ctx.name.compare("gfx900") == 0 || + ctx.name.compare("gfx903") == 0 || + ctx.name.compare("gfx905") == 0 + ) { /* Increase the number of threads for AMD VEGA gpus. * Limit the number of threads based on the issue: https://github.com/fireice-uk/xmr-stak/issues/5#issuecomment-339425089 @@ -109,6 +117,9 @@ private: */ maxThreads = 2024u; } + // increase all intensity limits by two for aeon + if(!::jconf::inst()->IsCurrencyMonero()) + maxThreads *= 2u; // keep 128MiB memory free (value is randomly chosen) size_t availableMem = ctx.freeMem - (128u * byteToMiB); @@ -118,14 +129,27 @@ private: size_t possibleIntensity = std::min( maxThreads , maxIntensity ); // map intensity to a multiple of the compute unit count, 8 is the number of threads per work group size_t intensity = (possibleIntensity / (8 * ctx.computeUnits)) * ctx.computeUnits * 8; - conf += std::string(" // gpu: ") + ctx.name + " memory:" + std::to_string(availableMem / byteToMiB) + "\n"; - conf += std::string(" // compute units: ") + std::to_string(ctx.computeUnits) + "\n"; - // set 8 threads per block (this is a good value for the most gpus) - conf += std::string(" { \"index\" : ") + std::to_string(ctx.deviceIdx) + ",\n" + - " \"intensity\" : " + std::to_string(intensity) + ", \"worksize\" : " + std::to_string(8) + ",\n" + - " \"affine_to_cpu\" : false, \"strided_index\" : true\n" - " },\n"; - ++i; + //If the intensity is 0, then it's because the multiple of the unit count is greater than intensity + if (intensity == 0) + { + printer::inst()->print_msg(L0, "WARNING: Auto detected intensity unexpectedly low. Try to set the environment variable GPU_SINGLE_ALLOC_PERCENT."); + intensity = possibleIntensity; + + } + if (intensity != 0) + { + conf += std::string(" // gpu: ") + ctx.name + " memory:" + std::to_string(availableMem / byteToMiB) + "\n"; + conf += std::string(" // compute units: ") + std::to_string(ctx.computeUnits) + "\n"; + // set 8 threads per block (this is a good value for the most gpus) + conf += std::string(" { \"index\" : ") + std::to_string(ctx.deviceIdx) + ",\n" + + " \"intensity\" : " + std::to_string(intensity) + ", \"worksize\" : " + std::to_string(8) + ",\n" + + " \"affine_to_cpu\" : false, \"strided_index\" : true\n" + " },\n"; + } + else + { + printer::inst()->print_msg(L0, "WARNING: Ignore gpu %s, %s MiB free memory is not enough to suggest settings.", ctx.name.c_str(), std::to_string(availableMem / byteToMiB).c_str()); + } } configTpl.replace("PLATFORMINDEX",std::to_string(platformIndex)); diff --git a/xmrstak/backend/amd/config.tpl b/xmrstak/backend/amd/config.tpl index af662f8..25b75a1 100644 --- a/xmrstak/backend/amd/config.tpl +++ b/xmrstak/backend/amd/config.tpl @@ -12,6 +12,9 @@ R"===( * [ * { "index" : 0, "intensity" : 1000, "worksize" : 8, "affine_to_cpu" : false, "strided_index" : true }, * ], + * If you do not wish to mine with your AMD GPU(s) then use: + * "gpu_threads_conf" : + * null, */ "gpu_threads_conf" : [ diff --git a/xmrstak/backend/amd/jconf.cpp b/xmrstak/backend/amd/jconf.cpp index 07afb19..f126342 100644 --- a/xmrstak/backend/amd/jconf.cpp +++ b/xmrstak/backend/amd/jconf.cpp @@ -56,9 +56,10 @@ struct configVal { Type iType; }; -//Same order as in configEnum, as per comment above +// Same order as in configEnum, as per comment above +// kNullType means any type configVal oConfigValues[] = { - { aGpuThreadsConf, "gpu_threads_conf", kArrayType }, + { aGpuThreadsConf, "gpu_threads_conf", kNullType }, { iPlatformIdx, "platform_index", kNumberType } }; @@ -68,6 +69,8 @@ inline bool checkType(Type have, Type want) { if(want == have) return true; + else if(want == kNullType) + return true; else if(want == kTrueType && have == kFalseType) return true; else if(want == kFalseType && have == kTrueType) diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp index e83527c..422c28c 100644 --- a/xmrstak/backend/amd/minethd.cpp +++ b/xmrstak/backend/amd/minethd.cpp @@ -139,7 +139,7 @@ std::vector<iBackend*>* minethd::thread_starter(uint32_t threadOffset, miner_wor if(cfg.cpu_aff >= 0) { #if defined(__APPLE__) - printer::inst()->print_msg(L1, "WARNING on MacOS thread affinity is only advisory."); + printer::inst()->print_msg(L1, "WARNING on macOS thread affinity is only advisory."); #endif printer::inst()->print_msg(L1, "Starting AMD GPU thread %d, affinity: %d.", i, (int)cfg.cpu_aff); @@ -245,7 +245,7 @@ void minethd::work_main() if ( (*((uint64_t*)(bResult + 24))) < oWork.iTarget) executor::inst()->push_event(ex_event(job_result(oWork.sJobID, results[i], bResult, iThreadNo), oWork.iPoolId)); else - executor::inst()->push_event(ex_event("AMD Invalid Result", oWork.iPoolId)); + executor::inst()->push_event(ex_event("AMD Invalid Result", pGpuCtx->deviceIdx, oWork.iPoolId)); } iCount += pGpuCtx->rawIntensity; diff --git a/xmrstak/backend/cpu/autoAdjust.hpp b/xmrstak/backend/cpu/autoAdjust.hpp index 7bdb14e..db805ec 100644 --- a/xmrstak/backend/cpu/autoAdjust.hpp +++ b/xmrstak/backend/cpu/autoAdjust.hpp @@ -33,25 +33,21 @@ class autoAdjust { public: - size_t hashMemSize; - size_t halfHashMemSize; - - autoAdjust() + bool printConfig() { + size_t hashMemSizeKB; + size_t halfHashMemSizeKB; + if(::jconf::inst()->IsCurrencyMonero()) { - hashMemSize = MONERO_MEMORY; - halfHashMemSize = hashMemSize / 2u; + hashMemSizeKB = MONERO_MEMORY / 1024u; + halfHashMemSizeKB = hashMemSizeKB / 2u; } else { - hashMemSize = AEON_MEMORY; - halfHashMemSize = hashMemSize / 2u; + hashMemSizeKB = AEON_MEMORY / 1024u; + halfHashMemSizeKB = hashMemSizeKB / 2u; } - } - - bool printConfig() - { configEditor configTpl{}; @@ -63,9 +59,10 @@ public: std::string conf; - if(!detectL3Size() || L3KB_size < halfHashMemSize || L3KB_size > (halfHashMemSize * 100u)) + + if(!detectL3Size() || L3KB_size < halfHashMemSizeKB || L3KB_size > (halfHashMemSizeKB * 2048u)) { - if(L3KB_size < halfHashMemSize || L3KB_size > (halfHashMemSize * 100)) + if(L3KB_size < halfHashMemSizeKB || L3KB_size > (halfHashMemSizeKB * 2048)) 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"); @@ -88,7 +85,7 @@ public: if(L3KB_size <= 0) break; - double_mode = L3KB_size / hashMemSize > (int32_t)(corecnt-i); + double_mode = L3KB_size / hashMemSizeKB > (int32_t)(corecnt-i); conf += std::string(" { \"low_power_mode\" : "); conf += std::string(double_mode ? "true" : "false"); @@ -107,9 +104,9 @@ public: aff_id++; if(double_mode) - L3KB_size -= hashMemSize * 2u; + L3KB_size -= hashMemSizeKB * 2u; else - L3KB_size -= hashMemSize; + L3KB_size -= hashMemSizeKB; } } @@ -142,7 +139,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)) / halfHashMemSize; + (get_masked(cpu_info[1], 11, 0) + 1) * (cpu_info[2] + 1)) / 1024; return true; } diff --git a/xmrstak/backend/cpu/config.tpl b/xmrstak/backend/cpu/config.tpl index b21a22d..fc4acb9 100644 --- a/xmrstak/backend/cpu/config.tpl +++ b/xmrstak/backend/cpu/config.tpl @@ -24,6 +24,9 @@ R"===( * { "low_power_mode" : false, "no_prefetch" : true, "affine_to_cpu" : 0 }, * { "low_power_mode" : false, "no_prefetch" : true, "affine_to_cpu" : 1 }, * ], + * If you do not wish to mine with your CPU(s) then use: + * "cpu_threads_conf" : + * null, */ "cpu_threads_conf" : diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h index 9b6e1dc..e4ccbc3 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h +++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h @@ -317,10 +317,9 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c _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 & MASK], _MM_HINT_T0); + bx0 = cx; uint64_t hi, lo, cl, ch; cl = ((uint64_t*)&l0[idx0 & MASK])[0]; @@ -329,15 +328,15 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c lo = _umul128(idx0, cl, &hi); al0 += hi; - ah0 += lo; ((uint64_t*)&l0[idx0 & MASK])[0] = al0; + al0 ^= cl; + if(PREFETCH) + _mm_prefetch((const char*)&l0[al0 & MASK], _MM_HINT_T0); + ah0 += lo; ((uint64_t*)&l0[idx0 & MASK])[1] = ah0; ah0 ^= ch; - al0 ^= cl; - idx0 = al0; - if(PREFETCH) - _mm_prefetch((const char*)&l0[idx0 & MASK], _MM_HINT_T0); + idx0 = al0; } // Optim - 90% time boundary diff --git a/xmrstak/backend/cpu/crypto/cryptonight_common.cpp b/xmrstak/backend/cpu/crypto/cryptonight_common.cpp index 8b2207d..1026b04 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight_common.cpp +++ b/xmrstak/backend/cpu/crypto/cryptonight_common.cpp @@ -31,6 +31,7 @@ extern "C" #include "cryptonight.h" #include "cryptonight_aesni.h" #include "xmrstak/backend/cryptonight.hpp" +#include "xmrstak/misc/console.hpp" #include "xmrstak/jconf.hpp" #include <stdio.h> #include <stdlib.h> @@ -73,6 +74,8 @@ void do_skein_hash(const void* input, size_t len, char* output) { void (* const extra_hashes[4])(const void *, size_t, char *) = {do_blake_hash, do_groestl_hash, do_jh_hash, do_skein_hash}; #ifdef _WIN32 +#include "xmrstak/misc/uac.hpp" + BOOL bRebootDesirable = FALSE; //If VirtualAlloc fails, suggest a reboot BOOL AddPrivilege(TCHAR* pszPrivilege) @@ -176,13 +179,16 @@ size_t cryptonight_init(size_t use_fast_mem, size_t use_mlock, alloc_msg* msg) if(AddPrivilege(TEXT("SeLockMemoryPrivilege")) == 0) { + printer::inst()->print_msg(L0, "Elevating because we need to set up fast memory privileges."); + RequestElevation(); + if(AddLargePageRights()) { msg->warning = "Added SeLockMemoryPrivilege to the current account. You need to reboot for it to work"; bRebootDesirable = TRUE; } else - msg->warning = "Obtaning SeLockMemoryPrivilege failed."; + msg->warning = "Obtaining SeLockMemoryPrivilege failed."; return 0; } @@ -247,6 +253,9 @@ cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, al #elif defined(__FreeBSD__) ptr->long_state = (uint8_t*)mmap(0, hashMemSize, PROT_READ | PROT_WRITE, MAP_PRIVATE | MAP_ANONYMOUS | MAP_ALIGNED_SUPER | MAP_PREFAULT_READ, -1, 0); +#elif defined(__OpenBSD__) + ptr->long_state = (uint8_t*)mmap(0, hashMemSize, PROT_READ | PROT_WRITE, + MAP_PRIVATE | MAP_ANON, -1, 0); #else ptr->long_state = (uint8_t*)mmap(0, hashMemSize, PROT_READ | PROT_WRITE, MAP_PRIVATE | MAP_ANONYMOUS | MAP_HUGETLB | MAP_POPULATE, 0, 0); diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp index 143b66f..cef4f8e 100644 --- a/xmrstak/backend/cpu/minethd.cpp +++ b/xmrstak/backend/cpu/minethd.cpp @@ -73,7 +73,16 @@ namespace cpu bool minethd::thd_setaffinity(std::thread::native_handle_type h, uint64_t cpu_id) { #if defined(_WIN32) - return SetThreadAffinityMask(h, 1ULL << cpu_id) != 0; + // we can only pin up to 64 threads + if(cpu_id < 64) + { + return SetThreadAffinityMask(h, 1ULL << cpu_id) != 0; + } + else + { + printer::inst()->print_msg(L0, "WARNING: Windows supports only affinity up to 63."); + return false; + } #elif defined(__APPLE__) thread_port_t mach_thread; thread_affinity_policy_data_t policy = { static_cast<integer_t>(cpu_id) }; @@ -84,6 +93,8 @@ bool minethd::thd_setaffinity(std::thread::native_handle_type h, uint64_t cpu_id CPU_ZERO(&mn); CPU_SET(cpu_id, &mn); return pthread_setaffinity_np(h, sizeof(cpuset_t), &mn) == 0; +#elif defined(__OpenBSD__) + printer::inst()->print_msg(L0,"WARNING: thread pinning is not supported under OPENBSD."); #else cpu_set_t mn; CPU_ZERO(&mn); @@ -307,7 +318,7 @@ std::vector<iBackend*> minethd::thread_starter(uint32_t threadOffset, miner_work if(cfg.iCpuAff >= 0) { #if defined(__APPLE__) - printer::inst()->print_msg(L1, "WARNING on MacOS thread affinity is only advisory."); + printer::inst()->print_msg(L1, "WARNING on macOS thread affinity is only advisory."); #endif printer::inst()->print_msg(L1, "Starting %dx thread, affinity: %d.", cfg.iMultiway, (int)cfg.iCpuAff); @@ -437,12 +448,13 @@ void minethd::work_main() globalStates::inst().calc_start_nonce(result.iNonce, oWork.bNiceHash, nonce_chunk); } - *piNonce = ++result.iNonce; + *piNonce = result.iNonce; hash_fun(oWork.bWorkBlob, oWork.iWorkSize, result.bResult, ctx); if (*piHashVal < oWork.iTarget) executor::inst()->push_event(ex_event(result, oWork.iPoolId)); + result.iNonce++; std::this_thread::yield(); } @@ -626,7 +638,7 @@ void minethd::multiway_work_main(cn_hash_fun_multi hash_fun_multi) } for (size_t i = 0; i < N; i++) - *piNonce[i] = ++iNonce; + *piNonce[i] = iNonce++; hash_fun_multi(bWorkBlob, oWork.iWorkSize, bHashOut, ctx); @@ -634,7 +646,7 @@ void minethd::multiway_work_main(cn_hash_fun_multi hash_fun_multi) { if (*piHashVal[i] < oWork.iTarget) { - executor::inst()->push_event(ex_event(job_result(oWork.sJobID, iNonce - N + 1 + i, bHashOut + 32 * i, iThreadNo), oWork.iPoolId)); + executor::inst()->push_event(ex_event(job_result(oWork.sJobID, iNonce - N + i, bHashOut + 32 * i, iThreadNo), oWork.iPoolId)); } } diff --git a/xmrstak/backend/nvidia/config.tpl b/xmrstak/backend/nvidia/config.tpl index 5479172..f489956 100644 --- a/xmrstak/backend/nvidia/config.tpl +++ b/xmrstak/backend/nvidia/config.tpl @@ -26,6 +26,9 @@ R"===( * "affine_to_cpu" : false, "sync_mode" : 3, * }, * ], + * If you do not wish to mine with your nVidia GPU(s) then use: + * "gpu_threads_conf" : + * null, */ "gpu_threads_conf" : diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp index 5564596..9fd08fb 100644 --- a/xmrstak/backend/nvidia/minethd.cpp +++ b/xmrstak/backend/nvidia/minethd.cpp @@ -166,7 +166,7 @@ std::vector<iBackend*>* minethd::thread_starter(uint32_t threadOffset, miner_wor if(cfg.cpu_aff >= 0) { #if defined(__APPLE__) - printer::inst()->print_msg(L1, "WARNING on MacOS thread affinity is only advisory."); + printer::inst()->print_msg(L1, "WARNING on macOS thread affinity is only advisory."); #endif printer::inst()->print_msg(L1, "Starting NVIDIA GPU thread %d, affinity: %d.", i, (int)cfg.cpu_aff); @@ -287,7 +287,7 @@ void minethd::work_main() if ( (*((uint64_t*)(bResult + 24))) < oWork.iTarget) executor::inst()->push_event(ex_event(job_result(oWork.sJobID, foundNonce[i], bResult, iThreadNo), oWork.iPoolId)); else - executor::inst()->push_event(ex_event("NVIDIA Invalid Result", oWork.iPoolId)); + executor::inst()->push_event(ex_event("NVIDIA Invalid Result", ctx.device_id, oWork.iPoolId)); } iCount += h_per_round; diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu index 15a6f36..cc97274 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu @@ -74,24 +74,36 @@ __device__ __forceinline__ uint64_t cuda_mul128( uint64_t multiplier, uint64_t m template< typename T > __device__ __forceinline__ T loadGlobal64( T * const addr ) { +#if (__CUDA_ARCH__ < 700) T x; asm volatile( "ld.global.cg.u64 %0, [%1];" : "=l"( x ) : "l"( addr ) ); return x; +#else + return *addr; +#endif } template< typename T > __device__ __forceinline__ T loadGlobal32( T * const addr ) { +#if (__CUDA_ARCH__ < 700) T x; asm volatile( "ld.global.cg.u32 %0, [%1];" : "=r"( x ) : "l"( addr ) ); return x; +#else + return *addr; +#endif } template< typename T > __device__ __forceinline__ void storeGlobal32( T* addr, T const & val ) { +#if (__CUDA_ARCH__ < 700) asm volatile( "st.global.cg.u32 [%0], %1;" : : "l"( addr ), "r"( val ) ); +#else + *addr = val; +#endif } template<size_t ITERATIONS, uint32_t THREAD_SHIFT> diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu index d865e13..92259db 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu @@ -204,7 +204,13 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx) break; }; - CUDA_CHECK(ctx->device_id, cudaDeviceSetCacheConfig(cudaFuncCachePreferL1)); + const int gpuArch = ctx->device_arch[0] * 10 + ctx->device_arch[1]; + + /* Disable L1 cache for GPUs before Volta. + * L1 speed is increased and latency reduced with Volta. + */ + if(gpuArch < 70) + CUDA_CHECK(ctx->device_id, cudaDeviceSetCacheConfig(cudaFuncCachePreferL1)); size_t hashMemSize; if(::jconf::inst()->IsCurrencyMonero()) @@ -441,6 +447,12 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) maxMemUsage = size_t(1024u) * byteToMiB; } + if(props.multiProcessorCount <= 6) + { + // limit memory usage for low end devices to reduce the number of threads + maxMemUsage = size_t(1024u) * byteToMiB; + } + int* tmp; cudaError_t err; // a device must be selected to get the right memory usage later on diff --git a/xmrstak/cli/cli-miner.cpp b/xmrstak/cli/cli-miner.cpp index b84b783..7dc0f2f 100644 --- a/xmrstak/cli/cli-miner.cpp +++ b/xmrstak/cli/cli-miner.cpp @@ -86,6 +86,9 @@ void help() cout<<" --noNVIDIA disable the NVIDIA miner backend"<<endl; cout<<" --nvidia FILE NVIDIA backend miner config file"<<endl; #endif +#ifndef CONF_NO_HTTPD + cout<<" -i --httpd HTTP_PORT HTTP interface port"<<endl; +#endif cout<<" "<<endl; cout<<"The following options can be used for automatic start without a guided config,"<<endl; cout<<"If config exists then this pool will be top priority."<<endl; @@ -209,6 +212,28 @@ void do_guided_config() currency = tmp; } + auto& http_port = params::inst().httpd_port; + if(http_port == params::httpd_port_unset) + { +#if defined(CONF_NO_HTTPD) + http_port = params::httpd_port_disabled; +#else + std::cout<<"- Do you want to use the HTTP interface?" <<std::endl; + std::cout<<"Unlike the screen display, browser interface is not affected by the GPU lag." <<std::endl; + std::cout<<"If you don't want to use it, please enter 0, otherwise enter port number that the miner should listen on" <<std::endl; + + int32_t port; + while(!(std::cin >> port) || port < 0 || port > 65535) + { + std::cin.clear(); + std::cin.ignore(INT_MAX, '\n'); + std::cout << "Invalid port number. Please enter a number between 0 and 65535." << std::endl; + } + + http_port = port; +#endif + } + auto& pool = params::inst().poolURL; bool userSetPool = true; if(pool.empty()) @@ -306,6 +331,7 @@ void do_guided_config() configTpl.replace("POOLCONF", pool_table); configTpl.replace("CURRENCY", currency); + configTpl.replace("HTTP_PORT", std::to_string(http_port)); configTpl.write(params::inst().configFile); std::cout<<"Configuration stored in file '"<<params::inst().configFile<<"'"<<std::endl; } @@ -342,7 +368,14 @@ int main(int argc, char *argv[]) params::inst().executablePrefix += seperator; } - bool uacDialog = true; + params::inst().minerArg0 = argv[0]; + params::inst().minerArgs.reserve(argc * 16); + for(int i = 1; i < argc; i++) + { + params::inst().minerArgs += " "; + params::inst().minerArgs += argv[i]; + } + bool pool_url_set = false; for(size_t i = 1; i < argc-1; i++) { @@ -504,9 +537,31 @@ int main(int argc, char *argv[]) } params::inst().configFile = argv[i]; } + else if(opName.compare("-i") == 0 || opName.compare("--httpd") == 0) + { + ++i; + if( i >=argc ) + { + printer::inst()->print_msg(L0, "No argument for parameter '-i/--httpd' given"); + win_exit(); + return 1; + } + + char* endp = nullptr; + long int ret = strtol(argv[i], &endp, 10); + + if(endp == nullptr || ret < 0 || ret > 65535) + { + printer::inst()->print_msg(L0, "Argument for parameter '-i/--httpd' must be a number between 0 and 65535"); + win_exit(); + return 1; + } + + params::inst().httpd_port = ret; + } else if(opName.compare("--noUAC") == 0) { - uacDialog = false; + params::inst().allowUAC = false; } else { @@ -516,20 +571,6 @@ int main(int argc, char *argv[]) } } -#ifdef _WIN32 - if(uacDialog && !IsElevated()) - { - std::string minerArgs; - for(int i = 1; i < argc; i++) - { - minerArgs += " "; - minerArgs += argv[i]; - } - - SelfElevate(argv[0], minerArgs); - } -#endif - // check if we need a guided start if(!configEditor::file_exist(params::inst().configFile)) do_guided_config(); @@ -540,22 +581,35 @@ int main(int argc, char *argv[]) return 1; } +#ifdef _WIN32 + /* For Windows 7 and 8 request elevation at all times unless we are using slow memory */ + if(jconf::inst()->GetSlowMemSetting() != jconf::slow_mem_cfg::always_use && !IsWindows10OrNewer()) + { + printer::inst()->print_msg(L0, "Elevating due to Windows 7 or 8. You need Windows 10 to use fast memory without UAC elevation."); + RequestElevation(); + } +#endif + if (!BackendConnector::self_test()) { win_exit(); return 1; } -#ifndef CONF_NO_HTTPD - if(jconf::inst()->GetHttpdPort() != 0) + if(jconf::inst()->GetHttpdPort() != uint16_t(params::httpd_port_disabled)) { +#ifdef CONF_NO_HTTPD + printer::inst()->print_msg(L0, "HTTPD port is enabled but this binary was compiled without HTTP support!"); + win_exit(); + return 1; +#else if (!httpd::inst()->start_daemon()) { win_exit(); return 1; } - } #endif + } printer::inst()->print_str("-------------------------------------------------------------------\n"); printer::inst()->print_str(get_version_str_short().c_str()); diff --git a/xmrstak/cli/xmr-stak.manifest b/xmrstak/cli/xmr-stak.manifest new file mode 100644 index 0000000..ed65c97 --- /dev/null +++ b/xmrstak/cli/xmr-stak.manifest @@ -0,0 +1,34 @@ +<?xml version="1.0" encoding="UTF-8" standalone="yes"?> +<assembly xmlns="urn:schemas-microsoft-com:asm.v1" manifestVersion="1.0" xmlns:asmv3="urn:schemas-microsoft-com:asm.v3"> +<assemblyIdentity + version="1.0.0.0" + processorArchitecture="amd64" + name="xmr-stak" + type="win32" +/> + <description>XMR-Stak Monero Miner</description> + <trustInfo xmlns="urn:schemas-microsoft-com:asm.v3"> + <security> + <requestedPrivileges> + <requestedExecutionLevel + level="asInvoker" + uiAccess="false" + /> + </requestedPrivileges> + </security> + </trustInfo> + <compatibility xmlns="urn:schemas-microsoft-com:compatibility.v1"> + <application> + <!-- Windows 10 --> + <supportedOS Id="{8e0f7a12-bfb3-4fe8-b9a5-48fd50a15a9a}"/> + <!-- Windows 8.1 --> + <supportedOS Id="{1f676c76-80e1-4239-95bb-83d0f6d0da78}"/> + <!-- Windows Vista --> + <supportedOS Id="{e2011457-1546-43c5-a5fe-008deee3d3f0}"/> + <!-- Windows 7 --> + <supportedOS Id="{35138b9a-5d96-4fbd-8e2d-a2440225f93a}"/> + <!-- Windows 8 --> + <supportedOS Id="{4a2f28e3-53b9-4441-ba9c-d69d4a4a6e38}"/> + </application> + </compatibility> +</assembly> diff --git a/xmrstak/config.tpl b/xmrstak/config.tpl index ae97190..2c7bd41 100644 --- a/xmrstak/config.tpl +++ b/xmrstak/config.tpl @@ -159,7 +159,7 @@ POOLCONF], * * httpd_port - Port we should listen on. Default, 0, will switch off the server. */ -"httpd_port" : 0, +"httpd_port" : HTTP_PORT, /* * HTTP Authentication diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp index f279f52..a1db451 100644 --- a/xmrstak/jconf.cpp +++ b/xmrstak/jconf.cpp @@ -242,7 +242,10 @@ uint64_t jconf::GetAutohashTime() uint16_t jconf::GetHttpdPort() { - return prv->configValues[iHttpdPort]->GetUint(); + if(xmrstak::params::inst().httpd_port == xmrstak::params::httpd_port_unset) + return prv->configValues[iHttpdPort]->GetUint(); + else + return uint16_t(xmrstak::params::inst().httpd_port); } const char* jconf::GetHttpUsername() diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp index a3088a5..055739b 100644 --- a/xmrstak/misc/executor.cpp +++ b/xmrstak/misc/executor.cpp @@ -594,7 +594,7 @@ void executor::ex_main() break; case EV_GPU_RES_ERROR: - log_result_error(std::string(ev.oGpuError.error_str)); + log_result_error(std::string(ev.oGpuError.error_str + std::string(" GPU ID ") + std::to_string(ev.oGpuError.idx))); break; case EV_PERF_TICK: @@ -761,6 +761,7 @@ void executor::hashrate_report(std::string& out) else out.append(1, '\n'); + double fTotalCur[3] = { 0.0, 0.0, 0.0}; for (i = 0; i < nthd; i++) { double fHps[3]; @@ -775,10 +776,14 @@ void executor::hashrate_report(std::string& out) out.append(hps_format(fHps[0], num, sizeof(num))).append(" |"); out.append(hps_format(fHps[1], num, sizeof(num))).append(" |"); out.append(hps_format(fHps[2], num, sizeof(num))).append(1, ' '); - - fTotal[0] += fHps[0]; - fTotal[1] += fHps[1]; - fTotal[2] += fHps[2]; + + fTotal[0] += (std::isnormal(fHps[0])) ? fHps[0] : 0.0; + fTotal[1] += (std::isnormal(fHps[1])) ? fHps[1] : 0.0; + fTotal[2] += (std::isnormal(fHps[2])) ? fHps[2] : 0.0; + + fTotalCur[0] += (std::isnormal(fHps[0])) ? fHps[0] : 0.0; + fTotalCur[1] += (std::isnormal(fHps[1])) ? fHps[1] : 0.0; + fTotalCur[2] += (std::isnormal(fHps[2])) ? fHps[2] : 0.0; if((i & 0x1) == 1) //Odd i's out.append("|\n"); @@ -786,21 +791,25 @@ void executor::hashrate_report(std::string& out) if((i & 0x1) == 1) //We had odd number of threads out.append("|\n"); - - if(nthd != 1) - out.append("-----------------------------------------------------\n"); - else - out.append("---------------------------\n"); + + out.append("Totals (").append(name).append("): "); + out.append(hps_format(fTotalCur[0], num, sizeof(num))); + out.append(hps_format(fTotalCur[1], num, sizeof(num))); + out.append(hps_format(fTotalCur[2], num, sizeof(num))); + out.append(" H/s\n"); + + out.append("-----------------------------------------------------------------\n"); } } - out.append("Totals: "); + out.append("Totals (ALL): "); out.append(hps_format(fTotal[0], num, sizeof(num))); out.append(hps_format(fTotal[1], num, sizeof(num))); out.append(hps_format(fTotal[2], num, sizeof(num))); out.append(" H/s\nHighest: "); out.append(hps_format(fHighestHps, num, sizeof(num))); out.append(" H/s\n"); + out.append("-----------------------------------------------------------------\n"); } char* time_format(char* buf, size_t len, std::chrono::system_clock::time_point time) diff --git a/xmrstak/misc/executor.hpp b/xmrstak/misc/executor.hpp index c2caa39..fbaa265 100644 --- a/xmrstak/misc/executor.hpp +++ b/xmrstak/misc/executor.hpp @@ -177,7 +177,6 @@ private: iPoolCallTimes.clear(); tPoolConnTime = std::chrono::system_clock::now(); iPoolHashes = 0; - iPoolDiff = 0; } double fHighestHps = 0.0; diff --git a/xmrstak/misc/uac.cpp b/xmrstak/misc/uac.cpp new file mode 100644 index 0000000..ad9d394 --- /dev/null +++ b/xmrstak/misc/uac.cpp @@ -0,0 +1,77 @@ +#ifdef _WIN32 +#include "xmrstak/misc/console.hpp" +#include "xmrstak/params.hpp" + +#include <string> +#include <windows.h> + +BOOL IsElevated() +{ + BOOL fRet = FALSE; + HANDLE hToken = NULL; + if (OpenProcessToken(GetCurrentProcess(), TOKEN_QUERY, &hToken)) + { + TOKEN_ELEVATION Elevation; + DWORD cbSize = sizeof(TOKEN_ELEVATION); + if (GetTokenInformation(hToken, TokenElevation, &Elevation, sizeof(Elevation), &cbSize)) + fRet = Elevation.TokenIsElevated; + } + if (hToken) + CloseHandle(hToken); + return fRet; +} + +BOOL SelfElevate(const std::string& my_path, const std::string& params) +{ + if (IsElevated()) + return FALSE; + + SHELLEXECUTEINFO shExecInfo = { 0 }; + shExecInfo.cbSize = sizeof(SHELLEXECUTEINFO); + shExecInfo.fMask = SEE_MASK_NOCLOSEPROCESS; + shExecInfo.hwnd = NULL; + shExecInfo.lpVerb = "runas"; + shExecInfo.lpFile = my_path.c_str(); + shExecInfo.lpParameters = params.c_str(); + shExecInfo.lpDirectory = NULL; + shExecInfo.nShow = SW_SHOW; + shExecInfo.hInstApp = NULL; + + if (!ShellExecuteEx(&shExecInfo)) + return FALSE; + + // Loiter in the background to make scripting easier + printer::inst()->print_msg(L0, "This window has been opened because xmr-stak needed to run as administrator. It can be safely closed now."); + WaitForSingleObject(shExecInfo.hProcess, INFINITE); + std::exit(0); + + return TRUE; +} + +VOID RequestElevation() +{ + if(IsElevated()) + return; + + if(!xmrstak::params::inst().allowUAC) + { + printer::inst()->print_msg(L0, "The miner needs to run as administrator, but you passed --noUAC option. Please remove it or set use_slow_memory to always."); + win_exit(); + return; + } + + SelfElevate(xmrstak::params::inst().minerArg0, xmrstak::params::inst().minerArgs); +} + +BOOL IsWindows10OrNewer() +{ + OSVERSIONINFOEX osvi = { 0 }; + osvi.dwOSVersionInfoSize = sizeof(OSVERSIONINFOEX); + osvi.dwMajorVersion = 10; + osvi.dwMinorVersion = 0; + DWORDLONG dwlConditionMask = 0; + VER_SET_CONDITION(dwlConditionMask, VER_MAJORVERSION, VER_GREATER_EQUAL); + VER_SET_CONDITION(dwlConditionMask, VER_MINORVERSION, VER_GREATER_EQUAL); + return ::VerifyVersionInfo(&osvi, VER_MAJORVERSION | VER_MINORVERSION, dwlConditionMask); +} +#endif diff --git a/xmrstak/misc/uac.hpp b/xmrstak/misc/uac.hpp index 55c5f1a..33c79ae 100644 --- a/xmrstak/misc/uac.hpp +++ b/xmrstak/misc/uac.hpp @@ -1,51 +1,10 @@ #pragma once #ifdef _WIN32 -#include "xmrstak/misc/console.hpp" - #include <string> -#include <windows.h> - -BOOL IsElevated() -{ - BOOL fRet = FALSE; - HANDLE hToken = NULL; - if (OpenProcessToken(GetCurrentProcess(), TOKEN_QUERY, &hToken)) - { - TOKEN_ELEVATION Elevation; - DWORD cbSize = sizeof(TOKEN_ELEVATION); - if (GetTokenInformation(hToken, TokenElevation, &Elevation, sizeof(Elevation), &cbSize)) - fRet = Elevation.TokenIsElevated; - } - if (hToken) - CloseHandle(hToken); - return fRet; -} - -BOOL SelfElevate(const char* my_path, const std::string& params) -{ - if (IsElevated()) - return FALSE; - - SHELLEXECUTEINFO shExecInfo = { 0 }; - shExecInfo.cbSize = sizeof(SHELLEXECUTEINFO); - shExecInfo.fMask = SEE_MASK_NOCLOSEPROCESS; - shExecInfo.hwnd = NULL; - shExecInfo.lpVerb = "runas"; - shExecInfo.lpFile = my_path; - shExecInfo.lpParameters = params.c_str(); - shExecInfo.lpDirectory = NULL; - shExecInfo.nShow = SW_SHOW; - shExecInfo.hInstApp = NULL; - - if (!ShellExecuteEx(&shExecInfo)) - return FALSE; - - // Loiter in the background to make scripting easier - printer::inst()->print_msg(L0, "This window has been opened because xmr-stak needed to run as administrator. It can be safely closed now."); - WaitForSingleObject(shExecInfo.hProcess, INFINITE); - std::exit(0); - return TRUE; -} +BOOL IsElevated(); +BOOL SelfElevate(const std::string& my_path, const std::string& params); +VOID RequestElevation(); +BOOL IsWindows10OrNewer(); #endif diff --git a/xmrstak/net/msgstruct.hpp b/xmrstak/net/msgstruct.hpp index a5affc8..8c4bdbe 100644 --- a/xmrstak/net/msgstruct.hpp +++ b/xmrstak/net/msgstruct.hpp @@ -66,8 +66,9 @@ struct sock_err // Unlike socket errors, GPU errors are read-only strings struct gpu_res_err { + size_t idx; // GPU index const char* error_str; - gpu_res_err(const char* error_str) : error_str(error_str) {} + gpu_res_err(const char* error_str, size_t idx) : error_str(error_str), idx(idx) {} }; enum ex_event_name { EV_INVALID_VAL, EV_SOCK_READY, EV_SOCK_ERROR, EV_GPU_RES_ERROR, @@ -99,7 +100,7 @@ struct ex_event }; ex_event() { iName = EV_INVALID_VAL; iPoolId = 0;} - ex_event(const char* gpu_err, size_t id) : iName(EV_GPU_RES_ERROR), iPoolId(id), oGpuError(gpu_err) {} + ex_event(const char* gpu_err, size_t gpu_idx, size_t id) : iName(EV_GPU_RES_ERROR), iPoolId(id), oGpuError(gpu_err, gpu_idx) {} ex_event(std::string&& err, bool silent, size_t id) : iName(EV_SOCK_ERROR), iPoolId(id), oSocketError(std::move(err), silent) { } ex_event(job_result dat, size_t id) : iName(EV_MINER_HAVE_RESULT), iPoolId(id), oJobResult(dat) {} ex_event(pool_job dat, size_t id) : iName(EV_POOL_HAVE_JOB), iPoolId(id), oPoolJob(dat) {} diff --git a/xmrstak/net/socket.cpp b/xmrstak/net/socket.cpp index e19d1d4..89e9902 100644 --- a/xmrstak/net/socket.cpp +++ b/xmrstak/net/socket.cpp @@ -190,7 +190,7 @@ void tls_socket::print_error() if(jconf::inst()->TlsSecureAlgos()) pCallback->set_socket_error("Unknown TLS error. Secure TLS maybe unsupported, try setting tls_secure_algo to false."); else - pCallback->set_socket_error("Unknown TLS error."); + pCallback->set_socket_error("Unknown TLS error. You might be trying to connect to a non-TLS port."); } else pCallback->set_socket_error(buf, len); diff --git a/xmrstak/params.hpp b/xmrstak/params.hpp index bc32612..d0d6685 100644 --- a/xmrstak/params.hpp +++ b/xmrstak/params.hpp @@ -31,6 +31,10 @@ struct params std::string poolUsername; bool nicehashMode = false; + static constexpr int32_t httpd_port_unset = -1; + static constexpr int32_t httpd_port_disabled = 0; + int32_t httpd_port = httpd_port_unset; + std::string currency; std::string configFile; @@ -38,6 +42,10 @@ struct params std::string configFileNVIDIA; std::string configFileCPU; + bool allowUAC = true; + std::string minerArg0; + std::string minerArgs; + params() : binaryName("xmr-stak"), executablePrefix(""), |
