diff options
46 files changed, 1397 insertions, 389 deletions
diff --git a/.travis.yml b/.travis.yml index 0d2d51a..4d53d48 100644 --- a/.travis.yml +++ b/.travis.yml @@ -70,10 +70,10 @@ matrix: - CMAKE_C_COMPILER=gcc-7 - XMRSTAK_CMAKE_FLAGS="-DCUDA_ENABLE=OFF -DOpenCL_ENABLE=OFF" - - os: osx - compiler: gcc - env: - - XMRSTAK_CMAKE_FLAGS="-DCUDA_ENABLE=OFF -DOpenCL_ENABLE=OFF" +# - os: osx +# compiler: gcc +# env: +# - XMRSTAK_CMAKE_FLAGS="-DCUDA_ENABLE=OFF -DOpenCL_ENABLE=OFF" before_install: - . CI/checkPRBranch 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 @@ -6,6 +6,11 @@ XMR-Stak is a universal Stratum pool miner. This miner supports CPUs, AMD and NV ## 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"> +## Video setup guide on Windows + +[<img src="https://gist.githubusercontent.com/fireice-uk/3621b179d56f57a8ead6303d8e415cf6/raw/3215c7a8078a038efb14b8dd8fc2ba0a7731605c/vidguidetmb.jpg">](https://www.youtube.com/watch?v=m9XFoQvLH8Y) +###### Video by Crypto Sewer + ## Overview * [Features](#features) * [Supported altcoins](#supported-altcoins) @@ -21,15 +26,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 +67,7 @@ fireice-uk: psychocrypt: ``` -43NoJVEXo21hGZ6tDG6Z3g4qimiGdJPE6GRxAmiWwm26gwr62Lqo7zRiCJFSBmbkwTGNuuES9ES5TgaVHceuYc4Y75txCTU +45tcqnJMgd3VqeTznNotiNj4G9PQoK67TGRiHyj6EYSZ31NUbAfs9XdiU5squmZb717iHJLxZv3KfEw8jCYGL5wa19yrVCn ``` ## Release Checksums diff --git a/THIRD-PARTY-LICENSES b/THIRD-PARTY-LICENSES index 3e62013..d3202f4 100644 --- a/THIRD-PARTY-LICENSES +++ b/THIRD-PARTY-LICENSES @@ -22,3 +22,8 @@ License: MIT License and BSD License ------------------------------------------------------------------------- +Package: PicoSHA2 +Authors: okdshin +License: MIT License + +------------------------------------------------------------------------- @@ -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 b7104ac..729f4d2 100644 --- a/doc/compile_Linux.md +++ b/doc/compile_Linux.md @@ -8,7 +8,7 @@ ### 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/Runtime @@ -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..129596c 100644 --- a/doc/compile_Windows.md +++ b/doc/compile_Windows.md @@ -4,53 +4,54 @@ ### Preparation -- open a command line `cmd` -- run `mkdir C:\xmr-stak-dep` +- Open a command line (Windows key + r) and enter `cmd` +- Execute `mkdir C:\xmr-stak-dep` -### Visual Studio 2017 Community +### Visual Studio Community 2017 -- download VS2017 Community and install from [https://www.visualstudio.com/downloads/](https://www.visualstudio.com/downloads/) -- during the install chose the components +- Download and install [Visual Studio Community 2017](https://www.visualstudio.com/downloads/) +- During install choose following components: - `Desktop development with C++` (left side) - - `VC++ 2015.3 v140 toolset for desktop` (right side) + - `VC++ 2015.3 v140 toolset for desktop` (right side - **NOT** needed for CUDA 9 or AMD GPU) + - 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 -- download and install the latest version from [https://cmake.org/download/](https://cmake.org/download/) -- tested version: [cmake 3.9](https://cmake.org/files/v3.9/cmake-3.9.0-rc3-win64-x64.msi) -- during the install choose the option `Add CMake to the system PATH for all users` +- Download and install latest version from https://cmake.org/download/ +- Tested version: [cmake 3.9](https://cmake.org/files/v3.9/cmake-3.9.0-rc3-win64-x64.msi) +- During install choose option: `Add CMake to the system PATH for all users` -### Cuda 8.0+ (only needed to use NVIDIA GPUs) +### Cuda 8.0+ (only needed for NVIDIA GPUs) -- donwload 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 +- Download and install https://developer.nvidia.com/cuda-downloads +- For minimal install choose `Custom installation options` during the install and select + - CUDA/Development - CUDA/Visual Studio Integration (ignore the warning during the install that VS2017 is not supported) - CUDA/Runtime - Driver components -### AMD APP SDK 3.0 (only needed to use AMD GPUs) +### AMD APP SDK 3.0 (only needed for AMD GPUs) -- download and install the latest version from [http://developer.amd.com/amd-accelerated-parallel-processing-app-sdk/](http://developer.amd.com/amd-accelerated-parallel-processing-app-sdk/) +- Download and install the latest version from http://developer.amd.com/amd-accelerated-parallel-processing-app-sdk/ ### Dependencies OpenSSL/Hwloc and Microhttpd -- for CUDA 8*: - - download the version 1 of the precompiled binary from [https://github.com/fireice-uk/xmr-stak-dep/releases/download/v1/xmr-stak-dep.zip](https://github.com/fireice-uk/xmr-stak-dep/releases/download/v1/xmr-stak-dep.zip) - - version 1 of the pre-compiled dependencies is not compatible with Visual Studio Toolset v141 -- for CUDA 9 **and/or** AMD GPUs, CPU: - - download the version 2 of the precompiled binary from [https://github.com/fireice-uk/xmr-stak-dep/releases/download/v2/xmr-stak-dep.zip](https://github.com/fireice-uk/xmr-stak-dep/releases/download/v2/xmr-stak-dep.zip) - - version 2 of the pre-compiled dependencies is not compatible with Visual Studio Toolset v140 -- unzip all to `C:\xmr-stak-dep` +- For CUDA 8*: + - Download version 1 of the precompiled binary from https://github.com/fireice-uk/xmr-stak-dep/releases/download/v1/xmr-stak-dep.zip + - Version 1 of the pre-compiled dependencies is not compatible with Visual Studio Toolset v141 +- For CUDA 9* **and/or** AMD GPUs, CPU: + - Download version 2 of the precompiled binary from https://github.com/fireice-uk/xmr-stak-dep/releases/download/v2/xmr-stak-dep.zip + - Version 2 of the pre-compiled dependencies is not compatible with Visual Studio Toolset v140 +- Extract archive to `C:\xmr-stak-dep` ### Validate the Dependency Folder -- open a command line `cmd` -- run +- Open a command line (Windows key + r) and enter `cmd` +- Execute ``` cd c:\xmr-stak-dep tree . ``` -- the result should have the same structure +- You should see something like this: ``` C:\xmr-stak-dep>tree . Folder PATH listing for volume Windows @@ -75,29 +76,37 @@ ## Compile -- download and unzip `xmr-stak` -- open the command line terminal `cmd` -- `cd` to your unzipped source code directory -- execute the following commands (NOTE: path to VS2017 can be different) +- 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 `Source Code.zip` (e.g. to `C:\Users\USERNAME\xmr-stak-<version>`) +- Open a command line (Windows key + r) and enter `cmd` +- Go to extracted source code directory (e.g. `cd C:\Users\USERNAME\xmr-stak-<version>`) +- Execute the following commands (NOTE: path to Visual Studio Community 2017 can be different) ``` + # Execute next line 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 + ``` +- Sometimes Windows will change the directory to `C:\Users\USERNAME\source\` instead of `C:\Users\USERNAME\xmr-stak-<version>\`. If that's the case execute `cd C:\Users\USERNAME\xmr-stak-<version>` followed by: + ``` mkdir build + cd build + + set CMAKE_PREFIX_PATH=C:\xmr-stak-dep\hwloc;C:\xmr-stak-dep\libmicrohttpd;C:\xmr-stak-dep\openssl ``` - - for CUDA 8* - ``` - cmake -G "Visual Studio 15 2017 Win64" -T v140,host=x64 .. - ``` - - for CUDA 9 **and/or** AMD GPUs, CPU - ``` - cmake -G "Visual Studio 15 2017 Win64" -T v141,host=x64 .. - ``` + +### CMake + +- See [build options](https://github.com/fireice-uk/xmr-stak/blob/master/doc/compile.md#build-system) to enable or disable dependencies. +- For CUDA 8* execute: `cmake -G "Visual Studio 15 2017 Win64" -T v140,host=x64 ..` +- For CUDA 9* **and/or** AMD GPUs, CPU execute: `cmake -G "Visual Studio 15 2017 Win64" -T v141,host=x64 ..` +- Then execute ``` cmake --build . --config Release --target install + cd bin\Release + copy C:\xmr-stak-dep\openssl\bin\* . ``` - -\* Miner is also compiled for AMD GPUs (if the AMD APP SDK is installed) and CPUs. -CUDA 8 requires a downgrade to the old v140 tool chain. +- Miner is by default compiled for NVIDIA GPUs (if CUDA is installed), AMD GPUs (if the AMD APP SDK is installed) and CPUs. 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..79afa00 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -15,6 +15,7 @@ #include "xmrstak/backend/cryptonight.hpp" #include "xmrstak/jconf.hpp" +#include "xmrstak/picosha2/picosha2.hpp" #include <stdio.h> #include <string.h> @@ -25,8 +26,41 @@ #include <regex> #include <cassert> +#include <fstream> +#include <sstream> +#include <vector> +#include <string> +#include <iostream> + +#if defined _MSC_VER +#include <direct.h> +#elif defined __GNUC__ +#include <sys/types.h> +#include <sys/stat.h> +#endif + + + #ifdef _WIN32 #include <windows.h> +#include <Shlobj.h> + +static inline void create_directory(std::string dirname) +{ + _mkdir(dirname.data()); +} + +static inline std::string get_home() +{ + char path[MAX_PATH + 1]; + // get folder "appdata\local" + if (SHGetSpecialFolderPathA(HWND_DESKTOP, path, CSIDL_LOCAL_APPDATA, FALSE)) + { + return path; + } + else + return "."; +} static inline void port_sleep(size_t sec) { @@ -34,6 +68,22 @@ static inline void port_sleep(size_t sec) } #else #include <unistd.h> +#include <pwd.h> + +static inline void create_directory(std::string dirname) +{ + mkdir(dirname.data(), 0744); +} + +static inline std::string get_home() +{ + const char *home = "."; + + if ((home = getenv("HOME")) == nullptr) + home = getpwuid(getuid())->pw_dir; + + return home; +} static inline void port_sleep(size_t sec) { @@ -84,6 +134,7 @@ const char* err_to_str(cl_int ret) return "CL_MISALIGNED_SUB_BUFFER_OFFSET"; case CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST: return "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST"; +#ifdef CL_VERSION_1_2 case CL_COMPILE_PROGRAM_FAILURE: return "CL_COMPILE_PROGRAM_FAILURE"; case CL_LINKER_NOT_AVAILABLE: @@ -94,6 +145,7 @@ const char* err_to_str(cl_int ret) return "CL_DEVICE_PARTITION_FAILED"; case CL_KERNEL_ARG_INFO_NOT_AVAILABLE: return "CL_KERNEL_ARG_INFO_NOT_AVAILABLE"; +#endif case CL_INVALID_VALUE: return "CL_INVALID_VALUE"; case CL_INVALID_DEVICE_TYPE: @@ -164,6 +216,7 @@ const char* err_to_str(cl_int ret) return "CL_INVALID_GLOBAL_WORK_SIZE"; case CL_INVALID_PROPERTY: return "CL_INVALID_PROPERTY"; +#ifdef CL_VERSION_1_2 case CL_INVALID_IMAGE_DESCRIPTOR: return "CL_INVALID_IMAGE_DESCRIPTOR"; case CL_INVALID_COMPILER_OPTIONS: @@ -172,6 +225,7 @@ const char* err_to_str(cl_int ret) return "CL_INVALID_LINKER_OPTIONS"; case CL_INVALID_DEVICE_PARTITION_COUNT: return "CL_INVALID_DEVICE_PARTITION_COUNT"; +#endif #if defined(CL_VERSION_2_0) && !defined(CONF_ENFORCE_OpenCL_1_2) case CL_INVALID_PIPE_SIZE: return "CL_INVALID_PIPE_SIZE"; @@ -323,57 +377,157 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_ return ERR_OCL_API; } - ctx->Program = clCreateProgramWithSource(opencl_ctx, 1, (const char**)&source_code, NULL, &ret); - if(ret != CL_SUCCESS) + std::vector<char> devNameVec(1024); + if((ret = clGetDeviceInfo(ctx->DeviceID, CL_DEVICE_NAME, devNameVec.size(), devNameVec.data(), NULL)) != CL_SUCCESS) { - printer::inst()->print_msg(L1,"Error %s when calling clCreateProgramWithSource on the contents of cryptonight.cl", err_to_str(ret)); + printer::inst()->print_msg(L1,"WARNING: %s when calling clGetDeviceInfo to get CL_DEVICE_NAME for device %u.", err_to_str(ret),ctx->deviceIdx ); return ERR_OCL_API; } char options[256]; - snprintf(options, sizeof(options), - "-DITERATIONS=%d -DMASK=%d -DWORKSIZE=%llu -DSTRIDED_INDEX=%d", - hasIterations, threadMemMask, int_port(ctx->workSize), ctx->stridedIndex ? 1 : 0); - ret = clBuildProgram(ctx->Program, 1, &ctx->DeviceID, options, NULL, NULL); - if(ret != CL_SUCCESS) + snprintf(options, sizeof(options), + "-DITERATIONS=%d -DMASK=%d -DWORKSIZE=%llu -DSTRIDED_INDEX=%d -DMEM_CHUNK_EXPONENT=%d -DCOMP_MODE=%d", + hasIterations, threadMemMask, int_port(ctx->workSize), ctx->stridedIndex, int(1u<<ctx->memChunk), ctx->compMode ? 1 : 0); + + /* create a hash for the compile time cache + * used data: + * - source code + * - device name + * - compile paramater + */ + std::string src_str(source_code); + src_str += options; + src_str += devNameVec.data(); + std::string hash_hex_str; + picosha2::hash256_hex_string(src_str, hash_hex_str); + + std::string cache_file = get_home() + "/.openclcache/" + hash_hex_str + ".openclbin"; + std::ifstream clBinFile(cache_file, std::ofstream::in | std::ofstream::binary); + if(!clBinFile.good()) { - size_t len; - printer::inst()->print_msg(L1,"Error %s when calling clBuildProgram.", err_to_str(ret)); + printer::inst()->print_msg(L1,"WARNING: OpenCL device %u - OpenCL binary %s not found.",ctx->deviceIdx, cache_file.c_str()); + ctx->Program = clCreateProgramWithSource(opencl_ctx, 1, (const char**)&source_code, NULL, &ret); + if(ret != CL_SUCCESS) + { + printer::inst()->print_msg(L1,"Error %s when calling clCreateProgramWithSource on the OpenCL miner code", err_to_str(ret)); + return ERR_OCL_API; + } - if((ret = clGetProgramBuildInfo(ctx->Program, ctx->DeviceID, CL_PROGRAM_BUILD_LOG, 0, NULL, &len)) != CL_SUCCESS) + ret = clBuildProgram(ctx->Program, 1, &ctx->DeviceID, options, NULL, NULL); + if(ret != CL_SUCCESS) { - printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for length of build log output.", err_to_str(ret)); + size_t len; + printer::inst()->print_msg(L1,"Error %s when calling clBuildProgram.", err_to_str(ret)); + + if((ret = clGetProgramBuildInfo(ctx->Program, ctx->DeviceID, CL_PROGRAM_BUILD_LOG, 0, NULL, &len)) != CL_SUCCESS) + { + printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for length of build log output.", err_to_str(ret)); + return ERR_OCL_API; + } + + char* BuildLog = (char*)malloc(len + 1); + BuildLog[0] = '\0'; + + if((ret = clGetProgramBuildInfo(ctx->Program, ctx->DeviceID, CL_PROGRAM_BUILD_LOG, len, BuildLog, NULL)) != CL_SUCCESS) + { + free(BuildLog); + printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for build log.", err_to_str(ret)); + return ERR_OCL_API; + } + + printer::inst()->print_str("Build log:\n"); + std::cerr<<BuildLog<<std::endl; + + free(BuildLog); return ERR_OCL_API; } - char* BuildLog = (char*)malloc(len + 1); - BuildLog[0] = '\0'; + cl_uint num_devices; + clGetProgramInfo(ctx->Program, CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint), &num_devices,NULL); - if((ret = clGetProgramBuildInfo(ctx->Program, ctx->DeviceID, CL_PROGRAM_BUILD_LOG, len, BuildLog, NULL)) != CL_SUCCESS) + + std::vector<cl_device_id> devices_ids(num_devices); + clGetProgramInfo(ctx->Program, CL_PROGRAM_DEVICES, sizeof(cl_device_id)* devices_ids.size(), devices_ids.data(),NULL); + int dev_id = 0; + /* Search for the gpu within the program context. + * The id can be different to ctx->DeviceID. + */ + for(auto & ocl_device : devices_ids) { - free(BuildLog); - printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for build log.", err_to_str(ret)); + if(ocl_device == ctx->DeviceID) + break; + dev_id++; + } + + cl_build_status status; + do + { + if((ret = clGetProgramBuildInfo(ctx->Program, ctx->DeviceID, CL_PROGRAM_BUILD_STATUS, sizeof(cl_build_status), &status, NULL)) != CL_SUCCESS) + { + printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for status of build.", err_to_str(ret)); + return ERR_OCL_API; + } + port_sleep(1); + } + while(status == CL_BUILD_IN_PROGRESS); + + std::vector<size_t> binary_sizes(num_devices); + clGetProgramInfo (ctx->Program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t) * binary_sizes.size(), binary_sizes.data(), NULL); + + std::vector<char*> all_programs(num_devices); + std::vector<std::vector<char>> program_storage; + + int p_id = 0; + size_t mem_size = 0; + // create memory structure to query all OpenCL program binaries + for(auto & p : all_programs) + { + program_storage.emplace_back(std::vector<char>(binary_sizes[p_id])); + all_programs[p_id] = program_storage[p_id].data(); + mem_size += binary_sizes[p_id]; + p_id++; + } + + if( ret = clGetProgramInfo(ctx->Program, CL_PROGRAM_BINARIES, num_devices * sizeof(char*), all_programs.data(),NULL) != CL_SUCCESS) + { + printer::inst()->print_msg(L1,"Error %s when calling clGetProgramInfo.", err_to_str(ret)); return ERR_OCL_API; } - - printer::inst()->print_str("Build log:\n"); - std::cerr<<BuildLog<<std::endl; - free(BuildLog); - return ERR_OCL_API; + std::ofstream file_stream; + std::cout<<get_home() + "/.openclcache/" + hash_hex_str + ".openclbin"<<std::endl; + file_stream.open(cache_file, std::ofstream::out | std::ofstream::binary); + file_stream.write(all_programs[dev_id], binary_sizes[dev_id]); + file_stream.close(); + printer::inst()->print_msg(L1, "OpenCL device %u - OpenCL binary file stored in file %s.",ctx->deviceIdx, cache_file.c_str()); } - - cl_build_status status; - do + else { - if((ret = clGetProgramBuildInfo(ctx->Program, ctx->DeviceID, CL_PROGRAM_BUILD_STATUS, sizeof(cl_build_status), &status, NULL)) != CL_SUCCESS) + printer::inst()->print_msg(L1, "OpenCL device %u - Load OpenCL binary file %s",ctx->deviceIdx, cache_file.c_str()); + std::ostringstream ss; + ss << clBinFile.rdbuf(); + std::string s = ss.str(); + + size_t bin_size = s.size(); + auto data_ptr = s.data(); + + cl_int clStatus; + ctx->Program = clCreateProgramWithBinary( + opencl_ctx, 1, &ctx->DeviceID, &bin_size, + (const unsigned char **)&data_ptr, &clStatus, &ret + ); + if(ret != CL_SUCCESS) { - printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for status of build.", err_to_str(ret)); + printer::inst()->print_msg(L1,"Error %s when calling clCreateProgramWithBinary. Try to delete file %s", err_to_str(ret), cache_file.c_str()); + return ERR_OCL_API; + } + ret = clBuildProgram(ctx->Program, 1, &ctx->DeviceID, NULL, NULL, NULL); + if(ret != CL_SUCCESS) + { + printer::inst()->print_msg(L1,"Error %s when calling clBuildProgram. Try to delete file %s", err_to_str(ret), cache_file.c_str()); return ERR_OCL_API; } - port_sleep(1); } - while(status == CL_BUILD_IN_PROGRESS); const char *KernelNames[] = { "cn0", "cn1", "cn2", "Blake", "Groestl", "JH", "Skein" }; for(int i = 0; i < 7; ++i) @@ -487,7 +641,7 @@ std::vector<GpuContext> getAMDDevices(int index) printer::inst()->print_msg(L1,"WARNING: %s when calling clGetDeviceInfo to get the device vendor name for device %u.", err_to_str(clStatus), k); continue; } - + std::string devVendor(devVendorVec.data()); if( devVendor.find("Advanced Micro Devices") != std::string::npos || devVendor.find("AMD") != std::string::npos) { @@ -518,13 +672,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); } } @@ -549,6 +703,8 @@ int getAMDPlatformIdx() clStatus = clGetPlatformIDs(numPlatforms, platforms, NULL); int platformIndex = -1; + // Mesa OpenCL is the fallback if no AMD or Apple OpenCL is found + int mesaPlatform = -1; if(clStatus == CL_SUCCESS) { @@ -559,13 +715,29 @@ int getAMDPlatformIdx() clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, infoSize, platformNameVec.data(), NULL); std::string platformName(platformNameVec.data()); - if( platformName.find("Advanced Micro Devices") != std::string::npos || platformName.find("Apple") != std::string::npos) + if( platformName.find("Advanced Micro Devices") != std::string::npos || + platformName.find("Apple") != std::string::npos || + platformName.find("Mesa") != std::string::npos + ) { - platformIndex = i; + printer::inst()->print_msg(L0,"Found AMD platform index id = %i, name = %s",i , platformName.c_str()); - break; + if(platformName.find("Mesa") != std::string::npos) + mesaPlatform = i; + else + { + // exit if AMD or Apple platform is found + platformIndex = i; + break; + } } } + // fall back to Mesa OpenCL + if(platformIndex == -1 && mesaPlatform != -1) + { + printer::inst()->print_msg(L0,"No AMD platform found select Mesa as OpenCL platform"); + platformIndex = mesaPlatform; + } } else printer::inst()->print_msg(L1,"WARNING: %s when calling clGetPlatformIDs for platform information.", err_to_str(clStatus)); @@ -694,8 +866,18 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx) source_code = std::regex_replace(source_code, std::regex("XMRSTAK_INCLUDE_BLAKE256"), blake256CL); source_code = std::regex_replace(source_code, std::regex("XMRSTAK_INCLUDE_GROESTL256"), groestl256CL); + // create a directory for the OpenCL compile cache + create_directory(get_home() + "/.openclcache"); + for(int i = 0; i < num_gpus; ++i) { + if(ctx[i].stridedIndex == 2 && (ctx[i].rawIntensity % ctx[i].workSize) != 0) + { + size_t reduced_intensity = (ctx[i].rawIntensity / ctx[i].workSize) * ctx[i].workSize; + ctx[i].rawIntensity = reduced_intensity; + printer::inst()->print_msg(L0, "WARNING AMD: gpu %d intensity is not a multiple of 'worksize', auto reduce intensity to %d", ctx[i].deviceIdx, int(reduced_intensity)); + } + if((ret = InitOpenCLGpu(opencl_ctx, &ctx[i], source_code.c_str())) != ERR_SUCCESS) { return ret; @@ -866,10 +1048,15 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput) size_t g_intensity = ctx->rawIntensity; size_t w_size = ctx->workSize; - // round up to next multiple of w_size - size_t g_thd = ((g_intensity + w_size - 1u) / w_size) * w_size; - // number of global threads must be a multiple of the work group size (w_size) - assert(g_thd%w_size == 0); + size_t g_thd = g_intensity; + + if(ctx->compMode) + { + // round up to next multiple of w_size + size_t g_thd = ((g_intensity + w_size - 1u) / w_size) * w_size; + // number of global threads must be a multiple of the work group size (w_size) + assert(g_thd%w_size == 0); + } for(int i = 2; i < 6; ++i) { diff --git a/xmrstak/backend/amd/amd_gpu/gpu.hpp b/xmrstak/backend/amd/amd_gpu/gpu.hpp index c17bac1..8fb7168 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.hpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.hpp @@ -25,6 +25,8 @@ struct GpuContext size_t rawIntensity; size_t workSize; int stridedIndex; + int memChunk; + int compMode; /*Output vars*/ cl_device_id DeviceID; diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl index 255fcbb..9383b04 100644 --- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl +++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl @@ -14,6 +14,11 @@ R"===( * along with this program. If not, see <http://www.gnu.org/licenses/>. */ +/* For Mesa clover support */ +#ifdef cl_clang_storage_class_specifiers +# pragma OPENCL EXTENSION cl_clang_storage_class_specifiers : enable +#endif + #ifdef cl_amd_media_ops #pragma OPENCL EXTENSION cl_amd_media_ops : enable #else @@ -399,7 +404,7 @@ static const __constant uchar rcon[8] = { 0x8d, 0x01, 0x02, 0x04, 0x08, 0x10, 0x void AESExpandKey256(uint *keybuf) { //#pragma unroll 4 - for(uint c = 8, i = 1; c < 60; ++c) + for(uint c = 8, i = 1; c < 40; ++c) { // For 256-bit keys, an sbox permutation is done every other 4th uint generated, AND every 8th uint t = ((!(c & 7)) || ((c & 7) == 4)) ? SubWord(keybuf[c - 1]) : keybuf[c - 1]; @@ -411,21 +416,32 @@ void AESExpandKey256(uint *keybuf) } } +#define MEM_CHUNK (1<<MEM_CHUNK_EXPONENT) + #if(STRIDED_INDEX==0) # define IDX(x) (x) -#else +#elif(STRIDED_INDEX==1) # define IDX(x) ((x) * (Threads)) +#elif(STRIDED_INDEX==2) +# define IDX(x) (((x) % MEM_CHUNK) + ((x) / MEM_CHUNK) * WORKSIZE * MEM_CHUNK) #endif +inline ulong getIdx() +{ +#if(STRIDED_INDEX==0 || STRIDED_INDEX==1 || STRIDED_INDEX==2) + return get_global_id(0) - get_global_offset(0); +#endif +} + __attribute__((reqd_work_group_size(WORKSIZE, 8, 1))) __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states, ulong Threads) { ulong State[25]; - uint ExpandedKey1[256]; + uint ExpandedKey1[40]; __local uint AES0[256], AES1[256], AES2[256], AES3[256]; uint4 text; - const ulong gIdx = get_global_id(0) - get_global_offset(0); + const ulong gIdx = getIdx(); for(int i = get_local_id(1) * WORKSIZE + get_local_id(0); i < 256; @@ -439,16 +455,20 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul } barrier(CLK_LOCAL_MEM_FENCE); - + +#if(COMP_MODE==1) // do not use early return here if(gIdx < Threads) +#endif { states += 25 * gIdx; #if(STRIDED_INDEX==0) Scratchpad += gIdx * (ITERATIONS >> 2); -#else +#elif(STRIDED_INDEX==1) Scratchpad += gIdx; +#elif(STRIDED_INDEX==2) + Scratchpad += get_group_id(0) * (ITERATIONS >> 2) * WORKSIZE + MEM_CHUNK * get_local_id(0); #endif ((ulong8 *)State)[0] = vload8(0, input); @@ -470,9 +490,10 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul } mem_fence(CLK_GLOBAL_MEM_FENCE); - +#if(COMP_MODE==1) // do not use early return here if(gIdx < Threads) +#endif { #pragma unroll for(int i = 0; i < 25; ++i) states[i] = State[i]; @@ -486,9 +507,10 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul } mem_fence(CLK_LOCAL_MEM_FENCE); - +#if(COMP_MODE==1) // do not use early return here if(gIdx < Threads) +#endif { #pragma unroll 2 for(int i = 0; i < (ITERATIONS >> 5); ++i) @@ -509,7 +531,7 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre ulong a[2], b[2]; __local uint AES0[256], AES1[256], AES2[256], AES3[256]; - const ulong gIdx = get_global_id(0) - get_global_offset(0); + const ulong gIdx = getIdx(); for(int i = get_local_id(0); i < 256; i += WORKSIZE) { @@ -523,15 +545,18 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre barrier(CLK_LOCAL_MEM_FENCE); uint4 b_x; - +#if(COMP_MODE==1) // do not use early return here if(gIdx < Threads) +#endif { states += 25 * gIdx; #if(STRIDED_INDEX==0) Scratchpad += gIdx * (ITERATIONS >> 2); -#else +#elif(STRIDED_INDEX==1) Scratchpad += gIdx; +#elif(STRIDED_INDEX==2) + Scratchpad += get_group_id(0) * (ITERATIONS >> 2) * WORKSIZE + MEM_CHUNK * get_local_id(0); #endif a[0] = states[0] ^ states[4]; @@ -544,8 +569,10 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre mem_fence(CLK_LOCAL_MEM_FENCE); +#if(COMP_MODE==1) // do not use early return here if(gIdx < Threads) +#endif { #pragma unroll 8 for(int i = 0; i < ITERATIONS; ++i) @@ -578,11 +605,11 @@ __attribute__((reqd_work_group_size(WORKSIZE, 8, 1))) __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global uint *Branch0, __global uint *Branch1, __global uint *Branch2, __global uint *Branch3, ulong Threads) { __local uint AES0[256], AES1[256], AES2[256], AES3[256]; - uint ExpandedKey2[256]; + uint ExpandedKey2[40]; ulong State[25]; uint4 text; - const ulong gIdx = get_global_id(0) - get_global_offset(0); + const ulong gIdx = getIdx(); for(int i = get_local_id(1) * WORKSIZE + get_local_id(0); i < 256; @@ -597,14 +624,18 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u barrier(CLK_LOCAL_MEM_FENCE); +#if(COMP_MODE==1) // do not use early return here if(gIdx < Threads) +#endif { states += 25 * gIdx; #if(STRIDED_INDEX==0) Scratchpad += gIdx * (ITERATIONS >> 2); -#else +#elif(STRIDED_INDEX==1) Scratchpad += gIdx; +#elif(STRIDED_INDEX==2) + Scratchpad += get_group_id(0) * (ITERATIONS >> 2) * WORKSIZE + MEM_CHUNK * get_local_id(0); #endif #if defined(__Tahiti__) || defined(__Pitcairn__) @@ -624,15 +655,17 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u barrier(CLK_LOCAL_MEM_FENCE); +#if(COMP_MODE==1) // do not use early return here if(gIdx < Threads) +#endif { #pragma unroll 2 for(int i = 0; i < (ITERATIONS >> 5); ++i) { text ^= Scratchpad[IDX((i << 3) + get_local_id(1))]; - #pragma unroll + #pragma unroll 10 for(int j = 0; j < 10; ++j) text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]); } @@ -642,8 +675,10 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u barrier(CLK_GLOBAL_MEM_FENCE); +#if(COMP_MODE==1) // do not use early return here if(gIdx < Threads) +#endif { if(!get_local_id(1)) { @@ -653,21 +688,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 +729,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 +739,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 +767,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 +801,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 +857,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 +933,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..8950105 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,28 @@ 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\" : 1, \"mem_chunk\" : 2,\n" + " \"comp_mode\" : 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..28855f0 100644 --- a/xmrstak/backend/amd/config.tpl +++ b/xmrstak/backend/amd/config.tpl @@ -1,17 +1,29 @@ R"===( /* * GPU configuration. You should play around with intensity and worksize as the fastest settings will vary. - * index - GPU index number usually starts from 0 - * intensity - Number of parallel GPU threads (nothing to do with CPU threads) - * worksize - Number of local GPU threads (nothing to do with CPU threads) + * index - GPU index number usually starts from 0 + * intensity - Number of parallel GPU threads (nothing to do with CPU threads) + * worksize - Number of local GPU threads (nothing to do with CPU threads) * affine_to_cpu - This will affine the thread to a CPU. This can make a GPU miner play along nicer with a CPU miner. * strided_index - switch memory pattern used for the scratch pad memory - * true = use 16byte contiguous memory per thread, the next memory block has offset of intensity blocks - * false = use a contiguous block of memory per thread + * 2 = chunked memory, chunk size is controlled by 'mem_chunk' + * required: intensity must be a multiple of worksize + * 1 or true = use 16byte contiguous memory per thread, the next memory block has offset of intensity blocks + * 0 or false = use a contiguous block of memory per thread + * mem_chunk - range 0 to 18: set the number of elements (16byte) per chunk + * this value is only used if 'strided_index' == 2 + * element count is computed with the equation: 2 to the power of 'mem_chunk' e.g. 4 means a chunk of 16 elements(256byte) + * comp_mode - Compatibility enable/disable the automatic guard around compute kernel which allows + * to use a intensity which is not the multiple of the worksize. + * If you set false and the intensity is not multiple of the worksize the miner can crash: + * in this case set the intensity to a multiple of the worksize or activate comp_mode. * "gpu_threads_conf" : * [ - * { "index" : 0, "intensity" : 1000, "worksize" : 8, "affine_to_cpu" : false, "strided_index" : true }, + * { "index" : 0, "intensity" : 1000, "worksize" : 8, "affine_to_cpu" : false, "strided_index" : true, "mem_chunk" : 2, "comp_mode" : 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..93ba709 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) @@ -103,14 +106,17 @@ bool jconf::GetThreadConfig(size_t id, thd_cfg &cfg) if(!oThdConf.IsObject()) return false; - const Value *idx, *intensity, *w_size, *aff, *stridedIndex; + const Value *idx, *intensity, *w_size, *aff, *stridedIndex, *memChunk, *compMode; idx = GetObjectMember(oThdConf, "index"); intensity = GetObjectMember(oThdConf, "intensity"); w_size = GetObjectMember(oThdConf, "worksize"); aff = GetObjectMember(oThdConf, "affine_to_cpu"); stridedIndex = GetObjectMember(oThdConf, "strided_index"); + memChunk = GetObjectMember(oThdConf, "mem_chunk"); + compMode = GetObjectMember(oThdConf, "comp_mode"); - if(idx == nullptr || intensity == nullptr || w_size == nullptr || aff == nullptr || stridedIndex == nullptr) + if(idx == nullptr || intensity == nullptr || w_size == nullptr || aff == nullptr || memChunk == nullptr || + stridedIndex == nullptr || compMode == nullptr) return false; if(!idx->IsUint64() || !intensity->IsUint64() || !w_size->IsUint64()) @@ -119,13 +125,38 @@ bool jconf::GetThreadConfig(size_t id, thd_cfg &cfg) if(!aff->IsUint64() && !aff->IsBool()) return false; - if(!stridedIndex->IsBool()) + if(!stridedIndex->IsBool() && !stridedIndex->IsNumber()) + { + printer::inst()->print_msg(L0, "ERROR: strided_index must be a bool or a number"); + return false; + } + + if(stridedIndex->IsBool()) + cfg.stridedIndex = stridedIndex->GetBool() ? 1 : 0; + else + cfg.stridedIndex = (int)stridedIndex->GetInt64(); + + if(cfg.stridedIndex > 2) + { + printer::inst()->print_msg(L0, "ERROR: strided_index must be smaller than 2"); + return false; + } + + cfg.memChunk = (int)memChunk->GetInt64(); + + if(!idx->IsUint64() || cfg.memChunk > 18 ) + { + printer::inst()->print_msg(L0, "ERROR: mem_chunk must be smaller than 18"); + return false; + } + + if(!compMode->IsBool()) return false; cfg.index = idx->GetUint64(); - cfg.intensity = intensity->GetUint64(); cfg.w_size = w_size->GetUint64(); - cfg.stridedIndex = stridedIndex->GetBool(); + cfg.intensity = intensity->GetUint64(); + cfg.compMode = compMode->GetBool(); if(aff->IsNumber()) cfg.cpu_aff = aff->GetInt64(); diff --git a/xmrstak/backend/amd/jconf.hpp b/xmrstak/backend/amd/jconf.hpp index ee1882a..580b69f 100644 --- a/xmrstak/backend/amd/jconf.hpp +++ b/xmrstak/backend/amd/jconf.hpp @@ -26,7 +26,9 @@ public: size_t intensity; size_t w_size; long long cpu_aff; - bool stridedIndex; + int stridedIndex; + int memChunk; + bool compMode; }; size_t GetThreadCount(); diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp index e83527c..8dfbce5 100644 --- a/xmrstak/backend/amd/minethd.cpp +++ b/xmrstak/backend/amd/minethd.cpp @@ -97,6 +97,8 @@ bool minethd::init_gpus() vGpuData[i].rawIntensity = cfg.intensity; vGpuData[i].workSize = cfg.w_size; vGpuData[i].stridedIndex = cfg.stridedIndex; + vGpuData[i].memChunk = cfg.memChunk; + vGpuData[i].compMode = cfg.compMode; } return InitOpenCL(vGpuData.data(), n, jconf::inst()->GetPlatformIdx()) == ERR_SUCCESS; @@ -139,7 +141,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 +247,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..cb4b950 100644 --- a/xmrstak/backend/cpu/config.tpl +++ b/xmrstak/backend/cpu/config.tpl @@ -2,7 +2,7 @@ R"===( /* * Thread configuration for each thread. Make sure it matches the number above. * low_power_mode - This can either be a boolean (true or false), or a number between 1 to 5. When set to true, - this mode will double the cache usage, and double the single thread performance. It will + * this mode will double the cache usage, and double the single thread performance. It will * consume much less power (as less cores are working), but will max out at around 80-85% of * the maximum performance. When set to a number N greater than 1, this mode will increase the * cache usage and single thread performance by N times. @@ -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..867a998 100644 --- a/xmrstak/backend/nvidia/minethd.cpp +++ b/xmrstak/backend/nvidia/minethd.cpp @@ -80,14 +80,22 @@ minethd::minethd(miner_work& pWork, size_t iNo, const jconf::thd_cfg& cfg) ctx.syncMode = cfg.syncMode; this->affinity = cfg.cpu_aff; - std::unique_lock<std::mutex> lck(thd_aff_set); - std::future<void> order_guard = order_fix.get_future(); + std::future<void> numa_guard = numa_promise.get_future(); + thread_work_guard = thread_work_promise.get_future(); oWorkThd = std::thread(&minethd::work_main, this); - order_guard.wait(); + /* Wait until the gpu memory is initialized and numa cpu memory is pinned. + * The startup time is reduced if the memory is initialized in sequential order + * without concurrent threads (CUDA driver is less occupied). + */ + numa_guard.wait(); +} - if(affinity >= 0) //-1 means no affinity +void minethd::start_mining() +{ + thread_work_promise.set_value(); + if(this->affinity >= 0) //-1 means no affinity if(!cpu::minethd::thd_setaffinity(oWorkThd.native_handle(), affinity)) printer::inst()->print_msg(L1, "WARNING setting affinity failed."); } @@ -166,7 +174,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); @@ -179,6 +187,11 @@ std::vector<iBackend*>* minethd::thread_starter(uint32_t threadOffset, miner_wor } + for (i = 0; i < n; i++) + { + static_cast<minethd*>((*pvThreads)[i])->start_mining(); + } + return pvThreads; } @@ -208,10 +221,18 @@ void minethd::work_main() if(affinity >= 0) //-1 means no affinity bindMemoryToNUMANode(affinity); - order_fix.set_value(); - std::unique_lock<std::mutex> lck(thd_aff_set); - lck.release(); + if(cuda_get_deviceinfo(&ctx) != 0 || cryptonight_extra_cpu_init(&ctx) != 1) + { + printer::inst()->print_msg(L0, "Setup failed for GPU %d. Exitting.\n", (int)iThreadNo); + std::exit(0); + } + + // numa memory bind and gpu memory is initialized + numa_promise.set_value(); + std::this_thread::yield(); + // wait until all NVIDIA devices are initialized + thread_work_guard.wait(); uint64_t iCount = 0; cryptonight_ctx* cpu_ctx; @@ -221,12 +242,6 @@ void minethd::work_main() globalStates::inst().iConsumeCnt++; - if(cuda_get_deviceinfo(&ctx) != 0 || cryptonight_extra_cpu_init(&ctx) != 1) - { - 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"); while (bQuit == 0) @@ -287,7 +302,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/minethd.hpp b/xmrstak/backend/nvidia/minethd.hpp index d13c868..fcd24fa 100644 --- a/xmrstak/backend/nvidia/minethd.hpp +++ b/xmrstak/backend/nvidia/minethd.hpp @@ -32,7 +32,8 @@ private: typedef void (*cn_hash_fun)(const void*, size_t, void*, cryptonight_ctx*); minethd(miner_work& pWork, size_t iNo, const jconf::thd_cfg& cfg); - + void start_mining(); + void work_main(); void consume_work(); @@ -44,8 +45,11 @@ private: static miner_work oGlobalWork; miner_work oWork; - std::promise<void> order_fix; - std::mutex thd_aff_set; + std::promise<void> numa_promise; + std::promise<void> thread_work_promise; + + // block thread until all NVIDIA GPUs are initialized + std::future<void> thread_work_guard; std::thread oWorkThd; int64_t affinity; 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..9053844 100644 --- a/xmrstak/cli/cli-miner.cpp +++ b/xmrstak/cli/cli-miner.cpp @@ -86,12 +86,16 @@ 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; cout<<" -o, --url URL pool url and port, e.g. pool.usxmrpool.com:3333"<<endl; cout<<" -O, --tls-url URL TLS pool url and port, e.g. pool.usxmrpool.com:10443"<<endl; cout<<" -u, --user USERNAME pool user name or wallet address"<<endl; + cout<<" -r, --rigid RIGID rig identifier for pool-side statistics (needs pool support)"<<endl; cout<<" -p, --pass PASSWD pool password, in the most cases x or empty \"\""<<endl; cout<<" --use-nicehash the pool should run in nicehash mode"<<endl; cout<<" \n"<<endl; @@ -143,6 +147,11 @@ std::string get_multipool_entry(bool& final) std::cin.clear(); std::cin.ignore(INT_MAX,'\n'); std::cout<<"- Password (mostly empty or x):"<<std::endl; getline(std::cin, passwd); + + std::string rigid; + std::cin.clear(); std::cin.ignore(INT_MAX,'\n'); + std::cout<<"- Rig identifier for pool-side statistics (needs pool support). Can be empty:"<<std::endl; + getline(std::cin, rigid); #ifdef CONF_NO_TLS bool tls = false; @@ -162,9 +171,9 @@ std::string get_multipool_entry(bool& final) final = !read_yes_no("- Do you want to add another pool? (y/n)"); - return "\t{\"pool_address\" : \"" + pool +"\", \"wallet_address\" : \"" + userName + "\", \"pool_password\" : \"" + - passwd + "\", \"use_nicehash\" : " + bool_to_str(nicehash) + ", \"use_tls\" : " + bool_to_str(tls) + - ", \"tls_fingerprint\" : \"\", \"pool_weight\" : " + std::to_string(pool_weight) + " },\n"; + return "\t{\"pool_address\" : \"" + pool +"\", \"wallet_address\" : \"" + userName + "\", \"rig_id\" : \"" + rigid + + "\", \"pool_password\" : \"" + passwd + "\", \"use_nicehash\" : " + bool_to_str(nicehash) + ", \"use_tls\" : " + + bool_to_str(tls) + ", \"tls_fingerprint\" : \"\", \"pool_weight\" : " + std::to_string(pool_weight) + " },\n"; } inline void prompt_once(bool& prompted) @@ -209,6 +218,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()) @@ -243,6 +274,17 @@ void do_guided_config() getline(std::cin, passwd); } + auto& rigid = params::inst().poolRigid; + if(rigid.empty() && !params::inst().userSetRigid) + { + prompt_once(prompted); + + // clear everything from stdin to allow an empty rigid + std::cin.clear(); std::cin.ignore(INT_MAX,'\n'); + std::cout<<"- Rig identifier for pool-side statistics (needs pool support). Can be empty:"<<std::endl; + getline(std::cin, rigid); + } + bool tls; #ifdef CONF_NO_TLS tls = false; @@ -290,9 +332,9 @@ void do_guided_config() pool_weight = 1; std::string pool_table; - pool_table += "\t{\"pool_address\" : \"" + pool +"\", \"wallet_address\" : \"" + userName + "\", \"pool_password\" : \"" + - passwd + "\", \"use_nicehash\" : " + bool_to_str(nicehash) + ", \"use_tls\" : " + bool_to_str(tls) + - ", \"tls_fingerprint\" : \"\", \"pool_weight\" : " + std::to_string(pool_weight) + " },\n"; + pool_table += "\t{\"pool_address\" : \"" + pool +"\", \"wallet_address\" : \"" + userName + "\", \"rig_id\" : \"" + rigid + + "\", \"pool_password\" : \"" + passwd + "\", \"use_nicehash\" : " + bool_to_str(nicehash) + ", \"use_tls\" : " + + bool_to_str(tls) + ", \"tls_fingerprint\" : \"\", \"pool_weight\" : " + std::to_string(pool_weight) + " },\n"; if(multipool) { @@ -306,6 +348,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 +385,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++) { @@ -489,6 +539,26 @@ int main(int argc, char *argv[]) params::inst().userSetPwd = true; params::inst().poolPasswd = argv[i]; } + else if(opName.compare("-r") == 0 || opName.compare("--rigid") == 0) + { + if(!pool_url_set) + { + printer::inst()->print_msg(L0, "Pool address has to be set if you want to specify rigid."); + win_exit(); + return 1; + } + + ++i; + if( i >=argc ) + { + printer::inst()->print_msg(L0, "No argument for parameter '-r/--rigid' given"); + win_exit(); + return 1; + } + + params::inst().userSetRigid = true; + params::inst().poolRigid = argv[i]; + } else if(opName.compare("--use-nicehash") == 0) { params::inst().nicehashMode = true; @@ -504,9 +574,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 +608,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 +618,38 @@ 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(strlen(jconf::inst()->GetOutputFile()) != 0) + printer::inst()->open_logfile(jconf::inst()->GetOutputFile()); + 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()); @@ -581,9 +675,6 @@ int main(int argc, char *argv[]) else printer::inst()->print_msg(L0,"Start mining: AEON"); - if(strlen(jconf::inst()->GetOutputFile()) != 0) - printer::inst()->open_logfile(jconf::inst()->GetOutputFile()); - executor::inst()->ex_start(jconf::inst()->DaemonMode()); uint64_t lastTime = get_timestamp_ms(); 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..451ea7b 100644 --- a/xmrstak/config.tpl +++ b/xmrstak/config.tpl @@ -2,6 +2,7 @@ R"===( /* * pool_address - Pool address should be in the form "pool.supportxmr.com:3333". Only stratum pools are supported. * wallet_address - Your wallet, or pool login. + * rig_id - Rig identifier for pool-side statistics (needs pool support). * pool_password - Can be empty in most cases or "x". * use_nicehash - Limit the nonce to 3 bytes as required by nicehash. * use_tls - This option will make us connect using Transport Layer Security. @@ -159,7 +160,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..c9d3a20 100644 --- a/xmrstak/jconf.cpp +++ b/xmrstak/jconf.cpp @@ -129,12 +129,13 @@ bool jconf::GetPoolConfig(size_t id, pool_cfg& cfg) return false; typedef const Value* cval; - cval jaddr, jlogin, jpasswd, jnicehash, jtls, jtlsfp, jwt; + cval jaddr, jlogin, jrigid, jpasswd, jnicehash, jtls, jtlsfp, jwt; const Value& oThdConf = prv->configValues[aPoolList]->GetArray()[id]; /* We already checked presence and types */ jaddr = GetObjectMember(oThdConf, "pool_address"); jlogin = GetObjectMember(oThdConf, "wallet_address"); + jrigid = GetObjectMember(oThdConf, "rig_id"); jpasswd = GetObjectMember(oThdConf, "pool_password"); jnicehash = GetObjectMember(oThdConf, "use_nicehash"); jtls = GetObjectMember(oThdConf, "use_tls"); @@ -143,6 +144,7 @@ bool jconf::GetPoolConfig(size_t id, pool_cfg& cfg) cfg.sPoolAddr = jaddr->GetString(); cfg.sWalletAddr = jlogin->GetString(); + cfg.sRigId = jrigid->GetString(); cfg.sPasswd = jpasswd->GetString(); cfg.nicehash = jnicehash->GetBool(); cfg.tls = jtls->GetBool(); @@ -242,7 +244,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() @@ -417,8 +422,8 @@ bool jconf::parse_config(const char* sFilename) std::vector<size_t> pool_weights; pool_weights.reserve(pool_cnt); - const char* aPoolValues[] = { "pool_address", "wallet_address", "pool_password", "use_nicehash", "use_tls", "tls_fingerprint", "pool_weight" }; - Type poolValTypes[] = { kStringType, kStringType, kStringType, kTrueType, kTrueType, kStringType, kNumberType }; + const char* aPoolValues[] = { "pool_address", "wallet_address", "rig_id", "pool_password", "use_nicehash", "use_tls", "tls_fingerprint", "pool_weight" }; + Type poolValTypes[] = { kStringType, kStringType, kStringType, kStringType, kTrueType, kTrueType, kStringType, kNumberType }; constexpr size_t pvcnt = sizeof(aPoolValues)/sizeof(aPoolValues[0]); for(uint32_t i=0; i < pool_cnt; i++) diff --git a/xmrstak/jconf.hpp b/xmrstak/jconf.hpp index df1bf79..9a4e958 100644 --- a/xmrstak/jconf.hpp +++ b/xmrstak/jconf.hpp @@ -23,6 +23,7 @@ public: struct pool_cfg { const char* sPoolAddr; const char* sWalletAddr; + const char* sRigId; const char* sPasswd; bool nicehash; bool tls; diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp index a3088a5..c4ba26e 100644 --- a/xmrstak/misc/executor.cpp +++ b/xmrstak/misc/executor.cpp @@ -518,13 +518,14 @@ void executor::ex_main() already_have_cli_pool = true; const char* wallet = params.poolUsername.empty() ? cfg.sWalletAddr : params.poolUsername.c_str(); + const char* rigid = params.userSetRigid ? params.poolRigid.c_str() : cfg.sRigId; const char* pwd = params.userSetPwd ? params.poolPasswd.c_str() : cfg.sPasswd; bool nicehash = cfg.nicehash || params.nicehashMode; - pools.emplace_back(i+1, cfg.sPoolAddr, wallet, pwd, 9.9, false, params.poolUseTls, cfg.tls_fingerprint, nicehash); + pools.emplace_back(i+1, cfg.sPoolAddr, wallet, rigid, pwd, 9.9, false, params.poolUseTls, cfg.tls_fingerprint, nicehash); } else - pools.emplace_back(i+1, cfg.sPoolAddr, cfg.sWalletAddr, cfg.sPasswd, cfg.weight, false, cfg.tls, cfg.tls_fingerprint, cfg.nicehash); + pools.emplace_back(i+1, cfg.sPoolAddr, cfg.sWalletAddr, cfg.sRigId, cfg.sPasswd, cfg.weight, false, cfg.tls, cfg.tls_fingerprint, cfg.nicehash); } if(!xmrstak::params::inst().poolURL.empty() && !already_have_cli_pool) @@ -536,22 +537,22 @@ void executor::ex_main() win_exit(); } - pools.emplace_back(i+1, params.poolURL.c_str(), params.poolUsername.c_str(), params.poolPasswd.c_str(), 9.9, false, params.poolUseTls, "", params.nicehashMode); + pools.emplace_back(i+1, params.poolURL.c_str(), params.poolUsername.c_str(), params.poolRigid.c_str(), params.poolPasswd.c_str(), 9.9, false, params.poolUseTls, "", params.nicehashMode); } if(jconf::inst()->IsCurrencyMonero()) { if(dev_tls) - pools.emplace_front(0, "donate.xmr-stak.net:6666", "", "", 0.0, true, true, "", false); + pools.emplace_front(0, "donate.xmr-stak.net:6666", "", "", "", 0.0, true, true, "", false); else - pools.emplace_front(0, "donate.xmr-stak.net:3333", "", "", 0.0, true, false, "", false); + pools.emplace_front(0, "donate.xmr-stak.net:3333", "", "", "", 0.0, true, false, "", false); } else { if(dev_tls) - pools.emplace_front(0, "donate.xmr-stak.net:7777", "", "", 0.0, true, true, "", true); + pools.emplace_front(0, "donate.xmr-stak.net:7777", "", "", "", 0.0, true, true, "", true); else - pools.emplace_front(0, "donate.xmr-stak.net:4444", "", "", 0.0, true, false, "", true); + pools.emplace_front(0, "donate.xmr-stak.net:4444", "", "", "", 0.0, true, false, "", true); } ex_event ev; @@ -594,7 +595,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 +762,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 +777,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 +792,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/jpsock.cpp b/xmrstak/net/jpsock.cpp index 7ee09e7..9c413dc 100644 --- a/xmrstak/net/jpsock.cpp +++ b/xmrstak/net/jpsock.cpp @@ -92,8 +92,8 @@ struct jpsock::opq_json_val opq_json_val(const Value* val) : val(val) {} }; -jpsock::jpsock(size_t id, const char* sAddr, const char* sLogin, const char* sPassword, double pool_weight, bool dev_pool, bool tls, const char* tls_fp, bool nicehash) : - net_addr(sAddr), usr_login(sLogin), usr_pass(sPassword), tls_fp(tls_fp), pool_id(id), pool_weight(pool_weight), pool(dev_pool), nicehash(nicehash), +jpsock::jpsock(size_t id, const char* sAddr, const char* sLogin, const char* sRigId, const char* sPassword, double pool_weight, bool dev_pool, bool tls, const char* tls_fp, bool nicehash) : + net_addr(sAddr), usr_login(sLogin), usr_rigid(sRigId), usr_pass(sPassword), tls_fp(tls_fp), pool_id(id), pool_weight(pool_weight), pool(dev_pool), nicehash(nicehash), connect_time(0), connect_attempts(0), disconnect_time(0), quiet_close(false) { sock_init(); @@ -523,8 +523,8 @@ bool jpsock::cmd_login() { char cmd_buffer[1024]; - snprintf(cmd_buffer, sizeof(cmd_buffer), "{\"method\":\"login\",\"params\":{\"login\":\"%s\",\"pass\":\"%s\",\"agent\":\"%s\"},\"id\":1}\n", - usr_login.c_str(), usr_pass.c_str(), get_version_str().c_str()); + snprintf(cmd_buffer, sizeof(cmd_buffer), "{\"method\":\"login\",\"params\":{\"login\":\"%s\",\"pass\":\"%s\",\"rigid\":\"%s\",\"agent\":\"%s\"},\"id\":1}\n", + usr_login.c_str(), usr_pass.c_str(), usr_rigid.c_str(), get_version_str().c_str()); opq_json_val oResult(nullptr); diff --git a/xmrstak/net/jpsock.hpp b/xmrstak/net/jpsock.hpp index 9d276b7..d9e5542 100644 --- a/xmrstak/net/jpsock.hpp +++ b/xmrstak/net/jpsock.hpp @@ -27,7 +27,7 @@ class base_socket; class jpsock { public: - jpsock(size_t id, const char* sAddr, const char* sLogin, const char* sPassword, double pool_weight, bool dev_pool, bool tls, const char* tls_fp, bool nicehash); + jpsock(size_t id, const char* sAddr, const char* sLogin, const char* sRigId, const char* sPassword, double pool_weight, bool dev_pool, bool tls, const char* tls_fp, bool nicehash); ~jpsock(); bool connect(std::string& sConnectError); @@ -82,6 +82,7 @@ public: private: std::string net_addr; std::string usr_login; + std::string usr_rigid; std::string usr_pass; std::string tls_fp; 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..bed3427 100644 --- a/xmrstak/params.hpp +++ b/xmrstak/params.hpp @@ -28,9 +28,15 @@ struct params std::string poolURL; bool userSetPwd = false; std::string poolPasswd; + bool userSetRigid = false; + std::string poolRigid; 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 +44,10 @@ struct params std::string configFileNVIDIA; std::string configFileCPU; + bool allowUAC = true; + std::string minerArg0; + std::string minerArgs; + params() : binaryName("xmr-stak"), executablePrefix(""), diff --git a/xmrstak/picosha2/picosha2.hpp b/xmrstak/picosha2/picosha2.hpp new file mode 100644 index 0000000..b9daec6 --- /dev/null +++ b/xmrstak/picosha2/picosha2.hpp @@ -0,0 +1,375 @@ +/* +The MIT License (MIT) + +Copyright (C) 2017 okdshin + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#ifndef PICOSHA2_H +#define PICOSHA2_H +// picosha2:20140213 + +#ifndef PICOSHA2_BUFFER_SIZE_FOR_INPUT_ITERATOR +#define PICOSHA2_BUFFER_SIZE_FOR_INPUT_ITERATOR \ + 1048576 //=1024*1024: default is 1MB memory +#endif + +#include <algorithm> +#include <cassert> +#include <iterator> +#include <sstream> +#include <vector> + +namespace picosha2 { +typedef unsigned long word_t; +typedef unsigned char byte_t; + +static const size_t k_digest_size = 32; + +namespace detail { +inline byte_t mask_8bit(byte_t x) { return x & 0xff; } + +inline word_t mask_32bit(word_t x) { return x & 0xffffffff; } + +const word_t add_constant[64] = { + 0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, + 0x923f82a4, 0xab1c5ed5, 0xd807aa98, 0x12835b01, 0x243185be, 0x550c7dc3, + 0x72be5d74, 0x80deb1fe, 0x9bdc06a7, 0xc19bf174, 0xe49b69c1, 0xefbe4786, + 0x0fc19dc6, 0x240ca1cc, 0x2de92c6f, 0x4a7484aa, 0x5cb0a9dc, 0x76f988da, + 0x983e5152, 0xa831c66d, 0xb00327c8, 0xbf597fc7, 0xc6e00bf3, 0xd5a79147, + 0x06ca6351, 0x14292967, 0x27b70a85, 0x2e1b2138, 0x4d2c6dfc, 0x53380d13, + 0x650a7354, 0x766a0abb, 0x81c2c92e, 0x92722c85, 0xa2bfe8a1, 0xa81a664b, + 0xc24b8b70, 0xc76c51a3, 0xd192e819, 0xd6990624, 0xf40e3585, 0x106aa070, + 0x19a4c116, 0x1e376c08, 0x2748774c, 0x34b0bcb5, 0x391c0cb3, 0x4ed8aa4a, + 0x5b9cca4f, 0x682e6ff3, 0x748f82ee, 0x78a5636f, 0x84c87814, 0x8cc70208, + 0x90befffa, 0xa4506ceb, 0xbef9a3f7, 0xc67178f2}; + +const word_t initial_message_digest[8] = {0x6a09e667, 0xbb67ae85, 0x3c6ef372, + 0xa54ff53a, 0x510e527f, 0x9b05688c, + 0x1f83d9ab, 0x5be0cd19}; + +inline word_t ch(word_t x, word_t y, word_t z) { return (x & y) ^ ((~x) & z); } + +inline word_t maj(word_t x, word_t y, word_t z) { + return (x & y) ^ (x & z) ^ (y & z); +} + +inline word_t rotr(word_t x, std::size_t n) { + assert(n < 32); + return mask_32bit((x >> n) | (x << (32 - n))); +} + +inline word_t bsig0(word_t x) { return rotr(x, 2) ^ rotr(x, 13) ^ rotr(x, 22); } + +inline word_t bsig1(word_t x) { return rotr(x, 6) ^ rotr(x, 11) ^ rotr(x, 25); } + +inline word_t shr(word_t x, std::size_t n) { + assert(n < 32); + return x >> n; +} + +inline word_t ssig0(word_t x) { return rotr(x, 7) ^ rotr(x, 18) ^ shr(x, 3); } + +inline word_t ssig1(word_t x) { return rotr(x, 17) ^ rotr(x, 19) ^ shr(x, 10); } + +template <typename RaIter1, typename RaIter2> +void hash256_block(RaIter1 message_digest, RaIter2 first, RaIter2 last) { + assert(first + 64 == last); + static_cast<void>(last); // for avoiding unused-variable warning + word_t w[64]; + std::fill(w, w + 64, 0); + for (std::size_t i = 0; i < 16; ++i) { + w[i] = (static_cast<word_t>(mask_8bit(*(first + i * 4))) << 24) | + (static_cast<word_t>(mask_8bit(*(first + i * 4 + 1))) << 16) | + (static_cast<word_t>(mask_8bit(*(first + i * 4 + 2))) << 8) | + (static_cast<word_t>(mask_8bit(*(first + i * 4 + 3)))); + } + for (std::size_t i = 16; i < 64; ++i) { + w[i] = mask_32bit(ssig1(w[i - 2]) + w[i - 7] + ssig0(w[i - 15]) + + w[i - 16]); + } + + word_t a = *message_digest; + word_t b = *(message_digest + 1); + word_t c = *(message_digest + 2); + word_t d = *(message_digest + 3); + word_t e = *(message_digest + 4); + word_t f = *(message_digest + 5); + word_t g = *(message_digest + 6); + word_t h = *(message_digest + 7); + + for (std::size_t i = 0; i < 64; ++i) { + word_t temp1 = h + bsig1(e) + ch(e, f, g) + add_constant[i] + w[i]; + word_t temp2 = bsig0(a) + maj(a, b, c); + h = g; + g = f; + f = e; + e = mask_32bit(d + temp1); + d = c; + c = b; + b = a; + a = mask_32bit(temp1 + temp2); + } + *message_digest += a; + *(message_digest + 1) += b; + *(message_digest + 2) += c; + *(message_digest + 3) += d; + *(message_digest + 4) += e; + *(message_digest + 5) += f; + *(message_digest + 6) += g; + *(message_digest + 7) += h; + for (std::size_t i = 0; i < 8; ++i) { + *(message_digest + i) = mask_32bit(*(message_digest + i)); + } +} + +} // namespace detail + +template <typename InIter> +void output_hex(InIter first, InIter last, std::ostream& os) { + os.setf(std::ios::hex, std::ios::basefield); + while (first != last) { + os.width(2); + os.fill('0'); + os << static_cast<unsigned int>(*first); + ++first; + } + os.setf(std::ios::dec, std::ios::basefield); +} + +template <typename InIter> +void bytes_to_hex_string(InIter first, InIter last, std::string& hex_str) { + std::ostringstream oss; + output_hex(first, last, oss); + hex_str.assign(oss.str()); +} + +template <typename InContainer> +void bytes_to_hex_string(const InContainer& bytes, std::string& hex_str) { + bytes_to_hex_string(bytes.begin(), bytes.end(), hex_str); +} + +template <typename InIter> +std::string bytes_to_hex_string(InIter first, InIter last) { + std::string hex_str; + bytes_to_hex_string(first, last, hex_str); + return hex_str; +} + +template <typename InContainer> +std::string bytes_to_hex_string(const InContainer& bytes) { + std::string hex_str; + bytes_to_hex_string(bytes, hex_str); + return hex_str; +} + +class hash256_one_by_one { + public: + hash256_one_by_one() { init(); } + + void init() { + buffer_.clear(); + std::fill(data_length_digits_, data_length_digits_ + 4, 0); + std::copy(detail::initial_message_digest, + detail::initial_message_digest + 8, h_); + } + + template <typename RaIter> + void process(RaIter first, RaIter last) { + add_to_data_length(std::distance(first, last)); + std::copy(first, last, std::back_inserter(buffer_)); + std::size_t i = 0; + for (; i + 64 <= buffer_.size(); i += 64) { + detail::hash256_block(h_, buffer_.begin() + i, + buffer_.begin() + i + 64); + } + buffer_.erase(buffer_.begin(), buffer_.begin() + i); + } + + void finish() { + byte_t temp[64]; + std::fill(temp, temp + 64, 0); + std::size_t remains = buffer_.size(); + std::copy(buffer_.begin(), buffer_.end(), temp); + temp[remains] = 0x80; + + if (remains > 55) { + std::fill(temp + remains + 1, temp + 64, 0); + detail::hash256_block(h_, temp, temp + 64); + std::fill(temp, temp + 64 - 4, 0); + } else { + std::fill(temp + remains + 1, temp + 64 - 4, 0); + } + + write_data_bit_length(&(temp[56])); + detail::hash256_block(h_, temp, temp + 64); + } + + template <typename OutIter> + void get_hash_bytes(OutIter first, OutIter last) const { + for (const word_t* iter = h_; iter != h_ + 8; ++iter) { + for (std::size_t i = 0; i < 4 && first != last; ++i) { + *(first++) = detail::mask_8bit( + static_cast<byte_t>((*iter >> (24 - 8 * i)))); + } + } + } + + private: + void add_to_data_length(word_t n) { + word_t carry = 0; + data_length_digits_[0] += n; + for (std::size_t i = 0; i < 4; ++i) { + data_length_digits_[i] += carry; + if (data_length_digits_[i] >= 65536u) { + carry = data_length_digits_[i] >> 16; + data_length_digits_[i] &= 65535u; + } else { + break; + } + } + } + void write_data_bit_length(byte_t* begin) { + word_t data_bit_length_digits[4]; + std::copy(data_length_digits_, data_length_digits_ + 4, + data_bit_length_digits); + + // convert byte length to bit length (multiply 8 or shift 3 times left) + word_t carry = 0; + for (std::size_t i = 0; i < 4; ++i) { + word_t before_val = data_bit_length_digits[i]; + data_bit_length_digits[i] <<= 3; + data_bit_length_digits[i] |= carry; + data_bit_length_digits[i] &= 65535u; + carry = (before_val >> (16 - 3)) & 65535u; + } + + // write data_bit_length + for (int i = 3; i >= 0; --i) { + (*begin++) = static_cast<byte_t>(data_bit_length_digits[i] >> 8); + (*begin++) = static_cast<byte_t>(data_bit_length_digits[i]); + } + } + std::vector<byte_t> buffer_; + word_t data_length_digits_[4]; // as 64bit integer (16bit x 4 integer) + word_t h_[8]; +}; + +inline void get_hash_hex_string(const hash256_one_by_one& hasher, + std::string& hex_str) { + byte_t hash[k_digest_size]; + hasher.get_hash_bytes(hash, hash + k_digest_size); + return bytes_to_hex_string(hash, hash + k_digest_size, hex_str); +} + +inline std::string get_hash_hex_string(const hash256_one_by_one& hasher) { + std::string hex_str; + get_hash_hex_string(hasher, hex_str); + return hex_str; +} + +namespace impl { +template <typename RaIter, typename OutIter> +void hash256_impl(RaIter first, RaIter last, OutIter first2, OutIter last2, int, + std::random_access_iterator_tag) { + hash256_one_by_one hasher; + // hasher.init(); + hasher.process(first, last); + hasher.finish(); + hasher.get_hash_bytes(first2, last2); +} + +template <typename InputIter, typename OutIter> +void hash256_impl(InputIter first, InputIter last, OutIter first2, + OutIter last2, int buffer_size, std::input_iterator_tag) { + std::vector<byte_t> buffer(buffer_size); + hash256_one_by_one hasher; + // hasher.init(); + while (first != last) { + int size = buffer_size; + for (int i = 0; i != buffer_size; ++i, ++first) { + if (first == last) { + size = i; + break; + } + buffer[i] = *first; + } + hasher.process(buffer.begin(), buffer.begin() + size); + } + hasher.finish(); + hasher.get_hash_bytes(first2, last2); +} +} + +template <typename InIter, typename OutIter> +void hash256(InIter first, InIter last, OutIter first2, OutIter last2, + int buffer_size = PICOSHA2_BUFFER_SIZE_FOR_INPUT_ITERATOR) { + picosha2::impl::hash256_impl( + first, last, first2, last2, buffer_size, + typename std::iterator_traits<InIter>::iterator_category()); +} + +template <typename InIter, typename OutContainer> +void hash256(InIter first, InIter last, OutContainer& dst) { + hash256(first, last, dst.begin(), dst.end()); +} + +template <typename InContainer, typename OutIter> +void hash256(const InContainer& src, OutIter first, OutIter last) { + hash256(src.begin(), src.end(), first, last); +} + +template <typename InContainer, typename OutContainer> +void hash256(const InContainer& src, OutContainer& dst) { + hash256(src.begin(), src.end(), dst.begin(), dst.end()); +} + +template <typename InIter> +void hash256_hex_string(InIter first, InIter last, std::string& hex_str) { + byte_t hashed[k_digest_size]; + hash256(first, last, hashed, hashed + k_digest_size); + std::ostringstream oss; + output_hex(hashed, hashed + k_digest_size, oss); + hex_str.assign(oss.str()); +} + +template <typename InIter> +std::string hash256_hex_string(InIter first, InIter last) { + std::string hex_str; + hash256_hex_string(first, last, hex_str); + return hex_str; +} + +inline void hash256_hex_string(const std::string& src, std::string& hex_str) { + hash256_hex_string(src.begin(), src.end(), hex_str); +} + +template <typename InContainer> +void hash256_hex_string(const InContainer& src, std::string& hex_str) { + hash256_hex_string(src.begin(), src.end(), hex_str); +} + +template <typename InContainer> +std::string hash256_hex_string(const InContainer& src) { + return hash256_hex_string(src.begin(), src.end()); +} + +} // namespace picosha2 + +#endif // PICOSHA2_H
\ No newline at end of file |
