diff options
| author | xmr-stak-devs <email@example.com> | 2018-03-25 13:21:57 +0100 |
|---|---|---|
| committer | fireice-uk <fireice-uk@users.noreply.github.com> | 2018-03-25 13:28:40 +0100 |
| commit | 1e7911e653a267ffd71199cdf7afaf1cfed5bad0 (patch) | |
| tree | 29efb953d9851b352298369104bf777fa3bf34e1 | |
| parent | 5014bdda628f64ab780d02de371bac4997573d10 (diff) | |
| download | xmr-stak-1e7911e653a267ffd71199cdf7afaf1cfed5bad0.zip xmr-stak-1e7911e653a267ffd71199cdf7afaf1cfed5bad0.tar.gz | |
XMR-Stak 2.3.0 RC
Co-authored-by: psychocrypt <psychocryptHPC@gmail.com>
Co-authored-by: fireice-uk <fireice-uk@users.noreply.github.com>
Co-authored-by: Lee Clagett <code@leeclagett.com>
Co-authored-by: curie-kief <curie-kief@users.noreply.github.com>
36 files changed, 1963 insertions, 599 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt index 3b3c7eb..15a2684 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -36,10 +36,6 @@ if(NOT CMAKE_BUILD_TYPE) endif() set_property(CACHE CMAKE_BUILD_TYPE PROPERTY STRINGS "${BUILD_TYPE}") -set(XMR-STAK_CURRENCY "all" CACHE STRING "select miner currency") -set_property(CACHE XMR-STAK_CURRENCY PROPERTY STRINGS "all;monero;aeon") - - set(XMR-STAK_COMPILE "native" CACHE STRING "select CPU compute architecture") set_property(CACHE XMR-STAK_COMPILE PROPERTY STRINGS "native;generic") if(XMR-STAK_COMPILE STREQUAL "native") @@ -53,16 +49,6 @@ else() message(FATAL_ERROR "XMR-STAK_COMPILE is set to an unknown value '${XMR-STAK_COMPILE}'") endif() -if(XMR-STAK_CURRENCY STREQUAL "all") - message(STATUS "Set miner currency to 'monero' and 'aeon'") -elseif(XMR-STAK_CURRENCY STREQUAL "aeon") - message(STATUS "Set miner currency to 'aeon'") - add_definitions("-DCONF_NO_MONERO=1") -elseif(XMR-STAK_CURRENCY STREQUAL "monero") - message(STATUS "Set miner currency to 'monero'") - add_definitions("-DCONF_NO_AEON=1") -endif() - # option to add static libgcc and libstdc++ option(CMAKE_LINK_STATIC "link as much as possible libraries static" OFF) @@ -438,6 +424,14 @@ else() endif() +# add -Wall for debug builds with gcc +if(CMAKE_CXX_COMPILER_ID STREQUAL "GNU") + if(CMAKE_BUILD_TYPE STREQUAL "Debug") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall") + set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -Wall") + endif() +endif() + # activate static libgcc and libstdc++ linking if(CMAKE_LINK_STATIC) set(BUILD_SHARED_LIBRARIES OFF) @@ -38,14 +38,20 @@ XMR-Stak is a universal Stratum pool miner. This miner supports CPUs, AMD and NV ## Supported altcoins -Besides Monero, following coins can be mined using this miner: +Besides [Monero](https://getmonero.org), following coins can be mined using this miner: -- [Aeon](http://www.aeon.cash/) +- [Aeon](http://www.aeon.cash) +- [Edollar](https://edollar.cash) - [Electroneum](https://electroneum.com) +- [Graft](https://www.graft.network) - [Intense](https://intensecoin.com) +- [Karbo](https://karbo.io) - [Sumokoin](https://www.sumokoin.org) -For all coins, except Aeon, you can use Monero settings. +If your prefered coin is not listed, you can chose one of the following algorithms: + +- Cryptonight - 2 MiB scratchpad memory +- Cryptonight-light - 1 MiB scratchpad memory Please note, this list is not complete, and is not an endorsement. @@ -8,6 +8,9 @@ * [Illegal instruction (core dumped)](#illegal-instruction) * [Virus Protection Alert](#virus-protection-alert) * [Change Currency to Mine](#change-currency-to-mine) +* [How can I mine Monero](#how-can-i-mine-monero) +* [Why is Monero named monero2](why-is-monero-named-monero2) +* [Which currency must be chosen if my fork coin is not listed](#which-currency-must-be-chosen-if-my-fork-coin-is-not-listed) ## "Obtaining SeLockMemoryPrivilege failed." @@ -70,3 +73,16 @@ If your antivirus software flags **xmr-stak**, it will likely move it to its qua If the miner is compiled for Monero and Aeon than you can change - the value `currency` in the config *or* - start the miner with the [command line option](usage.md) `--currency monero` or `--currency aeon` + +## How can I mine Monero + +Set the value `currency` in `pools.txt` to `monero2`. + +## Why is Monero named monero2 + +To avoid configuration conflicts after the hard fork of Monero to the new POW with our old naming schema where all cryptonight currencies was selected by choosing `monero` as currency we decided to switch to the name `monero2`. + +## Which currency must be chosen if my fork coin is not listed + +If your coin you want to mine is not listed please check the documentation of the coin and try to find out if `cryptonight` or `cryptonight-lite` is the used algorithm. +Select one of these generic coin algorithms. diff --git a/doc/compile.md b/doc/compile.md index 771c9d1..984c013 100644 --- a/doc/compile.md +++ b/doc/compile.md @@ -47,7 +47,6 @@ After the configuration you need to compile the miner, follow the guide for your - there is no *http* interface available if option is disabled: `cmake .. -DMICROHTTPD_ENABLE=OFF` - `OpenSSL_ENABLE` allow to disable/enable the dependency *OpenSSL* - it is not possible to connect to a *https* secured pool if option is disabled: `cmake .. -DOpenSSL_ENABLE=OFF` -- `XMR-STAK_CURRENCY` - compile for Monero(XMR) or Aeon(AEON) usage only e.g. `cmake .. -DXMR-STAK_CURRENCY=monero` - `XMR-STAK_COMPILE` select the CPU compute architecture (default: native) - native means the miner binary can be used only on the system where it is compiled but will archive the highest hash rate - use `cmake .. -DXMR-STAK_COMPILE=generic` to run the miner on all CPU's with sse2 diff --git a/doc/tuning.md b/doc/tuning.md index 5125387..47ad0bb 100644 --- a/doc/tuning.md +++ b/doc/tuning.md @@ -1,6 +1,7 @@ # Tuning Guide ## Content Overview +* [Benchmark](#benchmark) * [Windows](#windows) * [NVIDIA Backend](#nvidia-backend) * [Choose Value for `threads` and `blocks`](#choose-value-for-threads-and-blocks) @@ -8,11 +9,18 @@ * [AMD Backend](#amd-backend) * [Choose `intensity` and `worksize`](#choose-intensity-and-worksize) * [Add more GPUs](#add-more-gpus) + * [disable comp_mode](#disable-comp_mode) + * [change the scratchpad memory pattern](change-the-scratchpad-memory-pattern) * [Increase Memory Pool](#increase-memory-pool) * [Scratchpad Indexing](#scratchpad-indexing) * [CPU Backend](#cpu-backend) * [Choose Value for `low_power_mode`](#choose-value-for-low_power_mode) +## Benchmark +To benchmark the miner speed there are two ways. + - Mine against a pool end press the key `h` after 30 sec to see the hash report. + - Start the miner with the cli option `--benchmark BLOCKVERSION`. The miner will not connect to any pool and performs a 60sec performance benchmark with all enabled back-ends. + ## Windows "Run As Administrator" prompt (UAC) confirmation is needed to use large pages on Windows 7. On Windows 10 it is only needed once to set up the account to use them. @@ -46,8 +54,12 @@ To add a new GPU you need to add a new config set to `gpu_threads_conf`. ``` "gpu_threads_conf" : [ - { "index" : 0, "threads" : 17, "blocks" : 60, "bfactor" : 0, "bsleep" : 0, "affine_to_cpu" : false}, - { "index" : 1, "threads" : 17, "blocks" : 60, "bfactor" : 0, "bsleep" : 0, "affine_to_cpu" : false}, + { "index" : 0, "threads" : 17, "blocks" : 60, "bfactor" : 0, "bsleep" : 0, + "affine_to_cpu" : false, "sync_mode" : 3, + }, + { "index" : 1, "threads" : 17, "blocks" : 60, "bfactor" : 0, "bsleep" : 0, + "affine_to_cpu" : false, "sync_mode" : 3, + }, ], ``` @@ -70,13 +82,26 @@ If you are unsure of either GPU or platform index value, you can use `clinfo` to ``` "gpu_threads_conf" : [ - { "index" : 0, "intensity" : 1000, "worksize" : 8, "affine_to_cpu" : false }, - { "index" : 1, "intensity" : 1000, "worksize" : 8, "affine_to_cpu" : false }, + { "index" : 0, "intensity" : 1000, "worksize" : 8, "affine_to_cpu" : false, + "strided_index" : true, "mem_chunk" : 2, "comp_mode" : true + }, + { "index" : 1, "intensity" : 1000, "worksize" : 8, "affine_to_cpu" : false, + "strided_index" : true, "mem_chunk" : 2, "comp_mode" : true + }, ], "platform_index" : 0, ``` +### disable comp_mode + +`comp_mode` means compatibility mode and removes some checks in compute kernel those takes care that the miner can be used on a wide range of AMD/OpenCL GPU devices. +To avoid miner crashes the `intensity` should be a multiple of `worksize` if `comp_mode` is `false`. + +### change the scratchpad memory pattern + +By changing `strided_index` to `2` the number of contiguous elements (a 16 byte) for one miner thread can be fine tuned with the option `mem_chunk`. + ### Increase Memory Pool By setting the following environment variables before the miner is started OpenCl allows the miner to more threads. @@ -84,9 +109,9 @@ This variables must be set each time before the miner is started else it could b ``` export GPU_FORCE_64BIT_PTR=1 -export GPU_MAX_HEAP_SIZE=99 -export GPU_MAX_ALLOC_PERCENT=99 -export GPU_SINGLE_ALLOC_PERCENT=99 +export GPU_MAX_HEAP_SIZE=100 +export GPU_MAX_ALLOC_PERCENT=100 +export GPU_SINGLE_ALLOC_PERCENT=100 ``` *Note:* Windows user must use `set` instead of `export` to define an environment variable. diff --git a/doc/usage.md b/doc/usage.md index a810469..1f1fb09 100644 --- a/doc/usage.md +++ b/doc/usage.md @@ -5,7 +5,7 @@ * [Usage on Windows](#usage-on-windows) * [Usage on Linux](#usage-on-linux) * [Command Line Options](#command-line-options) -* [HTML and JSON API report configuraton](#xx) +* [HTML and JSON API report configuraton](#html-and-json-api-report-configuraton) ## Configurations @@ -13,12 +13,13 @@ Before you started the miner the first time there are no config files available. Config files will be created at the first start. The number of files depends on the available backends. `config.txt` contains the common miner settings. +`pools.txt` contains the selected mining pools and currency to mine. `amd.txt`, `cpu.txt` and `nvidia.txt` contains miner backend specific settings and can be used for further tuning ([Tuning Guide](tuning.md)). ## Usage on Windows 1) Double click the `xmr-stak.exe` file -2) Fill in the pool url, username and password +2) Fill in the pool url settings, currency, username and password `set XMRSTAK_NOWAIT=1` disable the dialog `Press any key to exit.` for non UAC execution. diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index c45f211..7547083 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -306,21 +306,9 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_ return ERR_OCL_API; } - size_t hashMemSize; - int threadMemMask; - int hasIterations; - if(::jconf::inst()->IsCurrencyMonero()) - { - hashMemSize = MONERO_MEMORY; - threadMemMask = MONERO_MASK; - hasIterations = MONERO_ITER; - } - else - { - hashMemSize = AEON_MEMORY; - threadMemMask = AEON_MASK; - hasIterations = AEON_ITER; - } + size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo()); + int threadMemMask = cn_select_mask(::jconf::inst()->GetMiningAlgo()); + int hashIterations = cn_select_iter(::jconf::inst()->GetMiningAlgo()); size_t g_thd = ctx->rawIntensity; ctx->ExtraBuffers[0] = clCreateBuffer(opencl_ctx, CL_MEM_READ_WRITE, hashMemSize * g_thd, NULL, &ret); @@ -384,11 +372,13 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_ return ERR_OCL_API; } - char options[256]; - 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); + auto miner_algo = ::jconf::inst()->GetMiningAlgo(); + char options[512]; + snprintf(options, sizeof(options), + "-DITERATIONS=%d -DMASK=%d -DWORKSIZE=%llu -DSTRIDED_INDEX=%d -DMEM_CHUNK_EXPONENT=%d -DCOMP_MODE=%d -DMEMORY=%llu -DALGO=%d", + hashIterations, threadMemMask, int_port(ctx->workSize), ctx->stridedIndex, int(1u<<ctx->memChunk), ctx->compMode ? 1 : 0, + int_port(hashMemSize), int(miner_algo)); /* create a hash for the compile time cache * used data: * - source code @@ -529,8 +519,8 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_ } } - const char *KernelNames[] = { "cn0", "cn1", "cn2", "Blake", "Groestl", "JH", "Skein" }; - for(int i = 0; i < 7; ++i) + const char *KernelNames[] = { "cn0", "cn1", "cn2", "Blake", "Groestl", "JH", "Skein", "cn1_monero" }; + for(int i = 0; i < 8; ++i) { ctx->Kernels[i] = clCreateKernel(ctx->Program, KernelNames[i], &ret); if(ret != CL_SUCCESS) @@ -887,7 +877,7 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx) return ERR_SUCCESS; } -size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target) +size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target, xmrstak_algo miner_algo, uint32_t version) { cl_int ret; @@ -932,29 +922,65 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar return(ERR_OCL_API); } - // CN2 Kernel + if(miner_algo == cryptonight_heavy) + { + // version + if ((ret = clSetKernelArg(ctx->Kernels[0], 4, sizeof(cl_uint), &version)) != CL_SUCCESS) + { + printer::inst()->print_msg(L1, "Error %s when calling clSetKernelArg for kernel 0, argument 4.", err_to_str(ret)); + return ERR_OCL_API; + } + } + + // CN1 Kernel + + /// @todo only activate if currency is monero + int cn_kernel_offset = 0; + if(miner_algo == cryptonight_monero && version >= 7) + { + cn_kernel_offset = 6; + } // Scratchpads - if((ret = clSetKernelArg(ctx->Kernels[1], 0, sizeof(cl_mem), ctx->ExtraBuffers + 0)) != CL_SUCCESS) + if((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 0, sizeof(cl_mem), ctx->ExtraBuffers + 0)) != CL_SUCCESS) { printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 1, argument 0.", err_to_str(ret)); return ERR_OCL_API; } // States - if((ret = clSetKernelArg(ctx->Kernels[1], 1, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS) + if((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 1, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS) { printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 1, argument 1.", err_to_str(ret)); return ERR_OCL_API; } // Threads - if((ret = clSetKernelArg(ctx->Kernels[1], 2, sizeof(cl_ulong), &numThreads)) != CL_SUCCESS) + if((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 2, sizeof(cl_ulong), &numThreads)) != CL_SUCCESS) { printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 1, argument 2.", err_to_str(ret)); return(ERR_OCL_API); } + if(miner_algo == cryptonight_monero && version >= 7) + { + // Input + if ((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS) + { + printer::inst()->print_msg(L1, "Error %s when calling clSetKernelArg for kernel 1, arugment 4(input buffer).", err_to_str(ret)); + return ERR_OCL_API; + } + } + else if(miner_algo == cryptonight_heavy) + { + // version + if ((ret = clSetKernelArg(ctx->Kernels[1], 3, sizeof(cl_uint), &version)) != CL_SUCCESS) + { + printer::inst()->print_msg(L1, "Error %s when calling clSetKernelArg for kernel 1, argument 3 (version).", err_to_str(ret)); + return ERR_OCL_API; + } + } + // CN3 Kernel // Scratchpads if((ret = clSetKernelArg(ctx->Kernels[2], 0, sizeof(cl_mem), ctx->ExtraBuffers + 0)) != CL_SUCCESS) @@ -1005,6 +1031,16 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar return(ERR_OCL_API); } + if(miner_algo == cryptonight_heavy) + { + // version + if ((ret = clSetKernelArg(ctx->Kernels[2], 7, sizeof(cl_uint), &version)) != CL_SUCCESS) + { + printer::inst()->print_msg(L1, "Error %s when calling clSetKernelArg for kernel 2, argument 7.", err_to_str(ret)); + return ERR_OCL_API; + } + } + for(int i = 0; i < 4; ++i) { // States @@ -1039,7 +1075,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar return ERR_SUCCESS; } -size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput) +size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo, uint32_t version) { cl_int ret; cl_uint zero = 0; @@ -1092,7 +1128,13 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput) }*/ size_t tmpNonce = ctx->Nonce; - if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[1], 1, &tmpNonce, &g_thd, &w_size, 0, NULL, NULL)) != CL_SUCCESS) + /// @todo only activate if currency is monero + int cn_kernel_offset = 0; + if(miner_algo == cryptonight_monero && version >= 7) + { + cn_kernel_offset = 6; + } + if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[1 + cn_kernel_offset], 1, &tmpNonce, &g_thd, &w_size, 0, NULL, NULL)) != CL_SUCCESS) { printer::inst()->print_msg(L1,"Error %s when calling clEnqueueNDRangeKernel for kernel %d.", err_to_str(ret), 1); return ERR_OCL_API; diff --git a/xmrstak/backend/amd/amd_gpu/gpu.hpp b/xmrstak/backend/amd/amd_gpu/gpu.hpp index 8fb7168..a387b15 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.hpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.hpp @@ -1,6 +1,7 @@ #pragma once #include "xmrstak/misc/console.hpp" +#include "xmrstak/jconf.hpp" #if defined(__APPLE__) #include <OpenCL/cl.h> @@ -35,7 +36,7 @@ struct GpuContext cl_mem OutputBuffer; cl_mem ExtraBuffers[6]; cl_program Program; - cl_kernel Kernels[7]; + cl_kernel Kernels[8]; size_t freeMem; int computeUnits; std::string name; @@ -49,7 +50,7 @@ int getAMDPlatformIdx(); std::vector<GpuContext> getAMDDevices(int index); size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx); -size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target); -size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput); +size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target, xmrstak_algo miner_algo, uint32_t version); +size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo, uint32_t version); diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl index 9383b04..7a36357 100644 --- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl +++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl @@ -433,8 +433,18 @@ inline ulong getIdx() #endif } +inline uint4 mix_and_propagate(__local uint4 xin[8][WORKSIZE]) +{ + return xin[(get_local_id(1)) % 8][get_local_id(0)] ^ xin[(get_local_id(1) + 1) % 8][get_local_id(0)]; +} + __attribute__((reqd_work_group_size(WORKSIZE, 8, 1))) -__kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states, ulong Threads) +__kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states, ulong Threads +// cryptonight_heavy +#if (ALGO == 4) + , uint version +#endif +) { ulong State[25]; uint ExpandedKey1[40]; @@ -464,11 +474,11 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul states += 25 * gIdx; #if(STRIDED_INDEX==0) - Scratchpad += gIdx * (ITERATIONS >> 2); + Scratchpad += gIdx * (MEMORY >> 4); #elif(STRIDED_INDEX==1) Scratchpad += gIdx; #elif(STRIDED_INDEX==2) - Scratchpad += get_group_id(0) * (ITERATIONS >> 2) * WORKSIZE + MEM_CHUNK * get_local_id(0); + Scratchpad += get_group_id(0) * (MEMORY >> 4) * WORKSIZE + MEM_CHUNK * get_local_id(0); #endif ((ulong8 *)State)[0] = vload8(0, input); @@ -507,13 +517,41 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul } mem_fence(CLK_LOCAL_MEM_FENCE); + +// cryptonight_heavy +#if (ALGO == 4) + if(version >= 3) + { + __local uint4 xin[8][WORKSIZE]; + + /* Also left over threads performe this loop. + * The left over thread results will be ignored + */ + for(size_t i=0; i < 16; i++) + { + #pragma unroll + for(int j = 0; j < 10; ++j) + text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey1)[j]); + barrier(CLK_LOCAL_MEM_FENCE); + xin[get_local_id(1)][get_local_id(0)] = text; + barrier(CLK_LOCAL_MEM_FENCE); + text = mix_and_propagate(xin); + } + } +#endif + #if(COMP_MODE==1) // do not use early return here if(gIdx < Threads) #endif { + int iterations = MEMORY >> 7; +#if (ALGO == 4) + if(version < 3) + iterations >>= 1; +#endif #pragma unroll 2 - for(int i = 0; i < (ITERATIONS >> 5); ++i) + for(int i = 0; i < iterations; ++i) { #pragma unroll for(int j = 0; j < 10; ++j) @@ -525,8 +563,22 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul mem_fence(CLK_GLOBAL_MEM_FENCE); } +#define VARIANT1_1(p) \ + uint table = 0x75310U; \ + uint index = (((p).s2 >> 26) & 12) | (((p).s2 >> 23) & 2); \ + (p).s2 ^= ((table >> index) & 0x30U) << 24 + +#define VARIANT1_2(p) ((uint2 *)&(p))[0] ^= tweak1_2 + +#define VARIANT1_INIT() \ + tweak1_2 = as_uint2(input[4]); \ + tweak1_2.s0 >>= 24; \ + tweak1_2.s0 |= tweak1_2.s1 << 8; \ + tweak1_2.s1 = get_global_id(0); \ + tweak1_2 ^= as_uint2(states[24]) + __attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) -__kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Threads) +__kernel void cn1_monero(__global uint4 *Scratchpad, __global ulong *states, ulong Threads, __global ulong *input) { ulong a[2], b[2]; __local uint AES0[256], AES1[256], AES2[256], AES3[256]; @@ -544,6 +596,7 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre barrier(CLK_LOCAL_MEM_FENCE); + uint2 tweak1_2; uint4 b_x; #if(COMP_MODE==1) // do not use early return here @@ -552,11 +605,11 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre { states += 25 * gIdx; #if(STRIDED_INDEX==0) - Scratchpad += gIdx * (ITERATIONS >> 2); + Scratchpad += gIdx * (MEMORY >> 4); #elif(STRIDED_INDEX==1) Scratchpad += gIdx; #elif(STRIDED_INDEX==2) - Scratchpad += get_group_id(0) * (ITERATIONS >> 2) * WORKSIZE + MEM_CHUNK * get_local_id(0); + Scratchpad += get_group_id(0) * (MEMORY >> 4) * WORKSIZE + MEM_CHUNK * get_local_id(0); #endif a[0] = states[0] ^ states[4]; @@ -565,6 +618,7 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre b[1] = states[3] ^ states[7]; b_x = ((uint4 *)b)[0]; + VARIANT1_INIT(); } mem_fence(CLK_LOCAL_MEM_FENCE); @@ -581,9 +635,10 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre ((uint4 *)c)[0] = Scratchpad[IDX((a[0] & MASK) >> 4)]; ((uint4 *)c)[0] = AES_Round(AES0, AES1, AES2, AES3, ((uint4 *)c)[0], ((uint4 *)a)[0]); - //b_x ^= ((uint4 *)c)[0]; - Scratchpad[IDX((a[0] & MASK) >> 4)] = b_x ^ ((uint4 *)c)[0]; + b_x ^= ((uint4 *)c)[0]; + VARIANT1_1(b_x); + Scratchpad[IDX((a[0] & MASK) >> 4)] = b_x; uint4 tmp; tmp = Scratchpad[IDX((c[0] & MASK) >> 4)]; @@ -591,18 +646,129 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre a[1] += c[0] * as_ulong2(tmp).s0; a[0] += mul_hi(c[0], as_ulong2(tmp).s0); + VARIANT1_2(a[1]); Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0]; + VARIANT1_2(a[1]); + + ((uint4 *)a)[0] ^= tmp; + + b_x = ((uint4 *)c)[0]; + } + } + mem_fence(CLK_GLOBAL_MEM_FENCE); +} + +__attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) +__kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Threads +// cryptonight_heavy +#if (ALGO == 4) + , uint version +#endif +) +{ + ulong a[2], b[2]; + __local uint AES0[256], AES1[256], AES2[256], AES3[256]; + + const ulong gIdx = getIdx(); + + for(int i = get_local_id(0); i < 256; i += WORKSIZE) + { + const uint tmp = AES0_C[i]; + AES0[i] = tmp; + AES1[i] = rotate(tmp, 8U); + AES2[i] = rotate(tmp, 16U); + AES3[i] = rotate(tmp, 24U); + } + + 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 * (MEMORY >> 4); +#elif(STRIDED_INDEX==1) + Scratchpad += gIdx; +#elif(STRIDED_INDEX==2) + Scratchpad += get_group_id(0) * (MEMORY >> 4) * WORKSIZE + MEM_CHUNK * get_local_id(0); +#endif + + a[0] = states[0] ^ states[4]; + b[0] = states[2] ^ states[6]; + a[1] = states[1] ^ states[5]; + b[1] = states[3] ^ states[7]; + + b_x = ((uint4 *)b)[0]; + } + + mem_fence(CLK_LOCAL_MEM_FENCE); + +#if(COMP_MODE==1) + // do not use early return here + if(gIdx < Threads) +#endif + { + ulong idx0 = a[0]; + ulong mask = MASK; + + int iterations = ITERATIONS; +#if (ALGO == 4) + if(version < 3) + { + iterations <<= 1; + mask -= 0x200000; + } +#endif + #pragma unroll 8 + for(int i = 0; i < iterations; ++i) + { + ulong c[2]; + + ((uint4 *)c)[0] = Scratchpad[IDX((idx0 & mask) >> 4)]; + ((uint4 *)c)[0] = AES_Round(AES0, AES1, AES2, AES3, ((uint4 *)c)[0], ((uint4 *)a)[0]); + //b_x ^= ((uint4 *)c)[0]; + + Scratchpad[IDX((idx0 & mask) >> 4)] = b_x ^ ((uint4 *)c)[0]; + + uint4 tmp; + tmp = Scratchpad[IDX((c[0] & mask) >> 4)]; + + a[1] += c[0] * as_ulong2(tmp).s0; + a[0] += mul_hi(c[0], as_ulong2(tmp).s0); + + Scratchpad[IDX((c[0] & mask) >> 4)] = ((uint4 *)a)[0]; ((uint4 *)a)[0] ^= tmp; + idx0 = a[0]; b_x = ((uint4 *)c)[0]; +// cryptonight_heavy +#if (ALGO == 4) + if(version >= 3) + { + long n = *((__global long*)(Scratchpad + (IDX((idx0 & mask) >> 4)))); + int d = ((__global int*)(Scratchpad + (IDX((idx0 & mask) >> 4))))[2]; + long q = n / (d | 0x5); + *((__global long*)(Scratchpad + (IDX((idx0 & mask) >> 4)))) = n ^ q; + idx0 = d ^ q; + } +#endif } } mem_fence(CLK_GLOBAL_MEM_FENCE); } __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) +__kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global uint *Branch0, __global uint *Branch1, __global uint *Branch2, __global uint *Branch3, ulong Threads +// cryptonight_heavy +#if (ALGO == 4) + , uint version +#endif + ) { __local uint AES0[256], AES1[256], AES2[256], AES3[256]; uint ExpandedKey2[40]; @@ -631,11 +797,11 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u { states += 25 * gIdx; #if(STRIDED_INDEX==0) - Scratchpad += gIdx * (ITERATIONS >> 2); + Scratchpad += gIdx * (MEMORY >> 4); #elif(STRIDED_INDEX==1) Scratchpad += gIdx; #elif(STRIDED_INDEX==2) - Scratchpad += get_group_id(0) * (ITERATIONS >> 2) * WORKSIZE + MEM_CHUNK * get_local_id(0); + Scratchpad += get_group_id(0) * (MEMORY >> 4) * WORKSIZE + MEM_CHUNK * get_local_id(0); #endif #if defined(__Tahiti__) || defined(__Pitcairn__) @@ -655,13 +821,67 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u barrier(CLK_LOCAL_MEM_FENCE); +#if (ALGO == 4) + __local uint4 xin[8][WORKSIZE]; +#endif + #if(COMP_MODE==1) // do not use early return here if(gIdx < Threads) #endif { + int iterations = MEMORY >> 7; +#if (ALGO == 4) + if(version < 3) + { + iterations >>= 1; + #pragma unroll 2 + for(int i = 0; i < iterations; ++i) + { + text ^= Scratchpad[IDX((i << 3) + get_local_id(1))]; + + #pragma unroll 10 + for(int j = 0; j < 10; ++j) + text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]); + } + } + else + { + #pragma unroll 2 + for(int i = 0; i < iterations; ++i) + { + text ^= Scratchpad[IDX((i << 3) + get_local_id(1))]; + + #pragma unroll 10 + for(int j = 0; j < 10; ++j) + text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]); + + + barrier(CLK_LOCAL_MEM_FENCE); + xin[get_local_id(1)][get_local_id(0)] = text; + barrier(CLK_LOCAL_MEM_FENCE); + text = mix_and_propagate(xin); + } + + #pragma unroll 2 + for(int i = 0; i < iterations; ++i) + { + text ^= Scratchpad[IDX((i << 3) + get_local_id(1))]; + + #pragma unroll 10 + for(int j = 0; j < 10; ++j) + text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]); + + + barrier(CLK_LOCAL_MEM_FENCE); + xin[get_local_id(1)][get_local_id(0)] = text; + barrier(CLK_LOCAL_MEM_FENCE); + text = mix_and_propagate(xin); + } + } +#else #pragma unroll 2 - for(int i = 0; i < (ITERATIONS >> 5); ++i) + for(int i = 0; i < iterations; ++i) { text ^= Scratchpad[IDX((i << 3) + get_local_id(1))]; @@ -669,7 +889,34 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u for(int j = 0; j < 10; ++j) text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]); } +#endif + } + +// cryptonight_heavy +#if (ALGO == 4) + if(version >= 3) + { + /* Also left over threads performe this loop. + * The left over thread results will be ignored + */ + for(size_t i=0; i < 16; i++) + { + #pragma unroll + for(int j = 0; j < 10; ++j) + text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]); + barrier(CLK_LOCAL_MEM_FENCE); + xin[get_local_id(1)][get_local_id(0)] = text; + barrier(CLK_LOCAL_MEM_FENCE); + text = mix_and_propagate(xin); + } + } +#endif +#if(COMP_MODE==1) + // do not use early return here + if(gIdx < Threads) +#endif + { vstore2(as_ulong2(text), get_local_id(1) + 4, states); } diff --git a/xmrstak/backend/amd/autoAdjust.hpp b/xmrstak/backend/amd/autoAdjust.hpp index 8950105..ea057a0 100644 --- a/xmrstak/backend/amd/autoAdjust.hpp +++ b/xmrstak/backend/amd/autoAdjust.hpp @@ -83,15 +83,7 @@ private: constexpr size_t byteToMiB = 1024u * 1024u; - size_t hashMemSize; - if(::jconf::inst()->IsCurrencyMonero()) - { - hashMemSize = MONERO_MEMORY; - } - else - { - hashMemSize = AEON_MEMORY; - } + size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo()); std::string conf; for(auto& ctx : devVec) @@ -118,7 +110,7 @@ private: maxThreads = 2024u; } // increase all intensity limits by two for aeon - if(!::jconf::inst()->IsCurrencyMonero()) + if(::jconf::inst()->GetMiningAlgo() == cryptonight_lite) maxThreads *= 2u; // keep 128MiB memory free (value is randomly chosen) diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp index 8dfbce5..46a04d5 100644 --- a/xmrstak/backend/amd/minethd.cpp +++ b/xmrstak/backend/amd/minethd.cpp @@ -191,9 +191,20 @@ void minethd::work_main() uint64_t iCount = 0; cryptonight_ctx* cpu_ctx; cpu_ctx = cpu::minethd::minethd_alloc_ctx(); - cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, ::jconf::inst()->IsCurrencyMonero()); + auto miner_algo = ::jconf::inst()->GetMiningAlgo(); + cn_hash_fun hash_fun; + if(miner_algo == cryptonight_monero || miner_algo == cryptonight_heavy) + { + // start with cryptonight and switch later if fork version is reached + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight); + } + else + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); + globalStates::inst().iConsumeCnt++; + uint8_t version = 0; + while (bQuit == 0) { if (oWork.bStall) @@ -207,6 +218,16 @@ void minethd::work_main() std::this_thread::sleep_for(std::chrono::milliseconds(100)); consume_work(); + uint8_t new_version = oWork.getVersion(); + if(miner_algo == cryptonight_monero && version < 7 && new_version >= 7) + { + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_monero); + } + else if(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3) + { + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy); + } + version = new_version; continue; } @@ -215,7 +236,8 @@ void minethd::work_main() assert(sizeof(job_result::sJobID) == sizeof(pool_job::sJobID)); uint64_t target = oWork.iTarget; - XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target); + /// \todo add monero hard for version + XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target, miner_algo, version); if(oWork.bNiceHash) pGpuCtx->Nonce = *(uint32_t*)(oWork.bWorkBlob + 39); @@ -231,7 +253,7 @@ void minethd::work_main() cl_uint results[0x100]; memset(results,0,sizeof(cl_uint)*(0x100)); - XMRRunJob(pGpuCtx, results); + XMRRunJob(pGpuCtx, results, miner_algo, version); for(size_t i = 0; i < results[0xFF]; i++) { @@ -258,6 +280,16 @@ void minethd::work_main() } consume_work(); + uint8_t new_version = oWork.getVersion(); + if(miner_algo == cryptonight_monero && version < 7 && new_version >= 7) + { + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_monero); + } + else if(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3) + { + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy); + } + version = new_version; } } diff --git a/xmrstak/backend/cpu/autoAdjustHwloc.hpp b/xmrstak/backend/cpu/autoAdjustHwloc.hpp index ddeb89b..568abb5 100644 --- a/xmrstak/backend/cpu/autoAdjustHwloc.hpp +++ b/xmrstak/backend/cpu/autoAdjustHwloc.hpp @@ -28,16 +28,8 @@ public: autoAdjust() { - if(::jconf::inst()->IsCurrencyMonero()) - { - hashMemSize = MONERO_MEMORY; - halfHashMemSize = hashMemSize / 2u; - } - else - { - hashMemSize = AEON_MEMORY; - halfHashMemSize = hashMemSize / 2u; - } + hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo()); + halfHashMemSize = hashMemSize / 2u; } bool printConfig() diff --git a/xmrstak/backend/cpu/crypto/cryptonight.h b/xmrstak/backend/cpu/crypto/cryptonight.h index 631c39a..5c9a733 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight.h +++ b/xmrstak/backend/cpu/crypto/cryptonight.h @@ -7,8 +7,6 @@ extern "C" { #include <stddef.h> #include <inttypes.h> -#include "xmrstak/backend/cryptonight.hpp" - typedef struct { uint8_t hash_state[224]; // Need only 200, explicit align diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h index e4ccbc3..85373e8 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h +++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h @@ -16,6 +16,7 @@ #pragma once #include "cryptonight.h" +#include "xmrstak/backend/cryptonight.hpp" #include <memory.h> #include <stdio.h> @@ -148,7 +149,20 @@ static inline void soft_aes_round(__m128i key, __m128i* x0, __m128i* x1, __m128i *x7 = soft_aesenc(*x7, key); } -template<size_t MEM, bool SOFT_AES, bool PREFETCH> +inline void mix_and_propagate(__m128i& x0, __m128i& x1, __m128i& x2, __m128i& x3, __m128i& x4, __m128i& x5, __m128i& x6, __m128i& x7) +{ + __m128i tmp0 = x0; + x0 = _mm_xor_si128(x0, x1); + x1 = _mm_xor_si128(x1, x2); + x2 = _mm_xor_si128(x2, x3); + x3 = _mm_xor_si128(x3, x4); + x4 = _mm_xor_si128(x4, x5); + x5 = _mm_xor_si128(x5, x6); + x6 = _mm_xor_si128(x6, x7); + x7 = _mm_xor_si128(x7, tmp0); +} + +template<size_t MEM, bool SOFT_AES, bool PREFETCH, xmrstak_algo ALGO> void cn_explode_scratchpad(const __m128i* input, __m128i* output) { // This is more than we have registers, compiler will assign 2 keys on the stack @@ -166,6 +180,40 @@ void cn_explode_scratchpad(const __m128i* input, __m128i* output) xin6 = _mm_load_si128(input + 10); xin7 = _mm_load_si128(input + 11); + if(ALGO == cryptonight_heavy) + { + for(size_t i=0; i < 16; i++) + { + if(SOFT_AES) + { + soft_aes_round(k0, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + soft_aes_round(k1, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + soft_aes_round(k2, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + soft_aes_round(k3, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + soft_aes_round(k4, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + soft_aes_round(k5, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + soft_aes_round(k6, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + soft_aes_round(k7, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + soft_aes_round(k8, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + soft_aes_round(k9, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + } + else + { + aes_round(k0, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + aes_round(k1, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + aes_round(k2, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + aes_round(k3, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + aes_round(k4, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + aes_round(k5, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + aes_round(k6, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + aes_round(k7, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + aes_round(k8, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + aes_round(k9, &xin0, &xin1, &xin2, &xin3, &xin4, &xin5, &xin6, &xin7); + } + mix_and_propagate(xin0, xin1, xin2, xin3, xin4, xin5, xin6, xin7); + } + } + for (size_t i = 0; i < MEM / sizeof(__m128i); i += 8) { if(SOFT_AES) @@ -213,7 +261,7 @@ void cn_explode_scratchpad(const __m128i* input, __m128i* output) } } -template<size_t MEM, bool SOFT_AES, bool PREFETCH> +template<size_t MEM, bool SOFT_AES, bool PREFETCH, xmrstak_algo ALGO> void cn_implode_scratchpad(const __m128i* input, __m128i* output) { // This is more than we have registers, compiler will assign 2 keys on the stack @@ -275,6 +323,93 @@ void cn_implode_scratchpad(const __m128i* input, __m128i* output) aes_round(k8, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); aes_round(k9, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); } + + if(ALGO == cryptonight_heavy) + mix_and_propagate(xout0, xout1, xout2, xout3, xout4, xout5, xout6, xout7); + } + + if(ALGO == cryptonight_heavy) + { + for (size_t i = 0; i < MEM / sizeof(__m128i); i += 8) + { + if(PREFETCH) + _mm_prefetch((const char*)input + i + 0, _MM_HINT_NTA); + + xout0 = _mm_xor_si128(_mm_load_si128(input + i + 0), xout0); + xout1 = _mm_xor_si128(_mm_load_si128(input + i + 1), xout1); + xout2 = _mm_xor_si128(_mm_load_si128(input + i + 2), xout2); + xout3 = _mm_xor_si128(_mm_load_si128(input + i + 3), xout3); + + if(PREFETCH) + _mm_prefetch((const char*)input + i + 4, _MM_HINT_NTA); + + xout4 = _mm_xor_si128(_mm_load_si128(input + i + 4), xout4); + xout5 = _mm_xor_si128(_mm_load_si128(input + i + 5), xout5); + xout6 = _mm_xor_si128(_mm_load_si128(input + i + 6), xout6); + xout7 = _mm_xor_si128(_mm_load_si128(input + i + 7), xout7); + + if(SOFT_AES) + { + soft_aes_round(k0, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k1, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k2, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k3, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k4, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k5, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k6, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k7, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k8, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k9, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + } + else + { + aes_round(k0, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k1, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k2, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k3, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k4, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k5, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k6, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k7, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k8, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k9, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + } + + if(ALGO == cryptonight_heavy) + mix_and_propagate(xout0, xout1, xout2, xout3, xout4, xout5, xout6, xout7); + } + + for(size_t i=0; i < 16; i++) + { + if(SOFT_AES) + { + soft_aes_round(k0, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k1, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k2, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k3, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k4, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k5, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k6, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k7, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k8, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + soft_aes_round(k9, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + } + else + { + aes_round(k0, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k1, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k2, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k3, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k4, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k5, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k6, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k7, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k8, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + aes_round(k9, &xout0, &xout1, &xout2, &xout3, &xout4, &xout5, &xout6, &xout7); + } + + mix_and_propagate(xout0, xout1, xout2, xout3, xout4, xout5, xout6, xout7); + } } _mm_store_si128(output + 4, xout0); @@ -287,13 +422,45 @@ void cn_implode_scratchpad(const __m128i* input, __m128i* output) _mm_store_si128(output + 11, xout7); } -template<size_t MASK, size_t ITERATIONS, size_t MEM, bool SOFT_AES, bool PREFETCH> +inline void cryptonight_monero_tweak(uint64_t* mem_out, __m128i tmp) +{ + mem_out[0] = _mm_cvtsi128_si64(tmp); + + tmp = _mm_castps_si128(_mm_movehl_ps(_mm_castsi128_ps(tmp), _mm_castsi128_ps(tmp))); + uint64_t vh = _mm_cvtsi128_si64(tmp); + + uint8_t x = vh >> 24; + static const uint16_t table = 0x7531; + const uint8_t index = (((x >> 3) & 6) | (x & 1)) << 1; + vh ^= ((table >> index) & 0x3) << 28; + + mem_out[1] = vh; +} + +template<xmrstak_algo ALGO, bool SOFT_AES, bool PREFETCH> void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_ctx* ctx0) { + constexpr size_t MASK = cn_select_mask<ALGO>(); + constexpr size_t ITERATIONS = cn_select_iter<ALGO>(); + constexpr size_t MEM = cn_select_memory<ALGO>(); + + if(ALGO == cryptonight_monero && len < 43) + { + memset(output, 0, 32); + return; + } + keccak((const uint8_t *)input, len, ctx0->hash_state, 200); + uint64_t monero_const; + if(ALGO == cryptonight_monero) + { + monero_const = *reinterpret_cast<const uint64_t*>(reinterpret_cast<const uint8_t*>(input) + 35); + monero_const ^= *(reinterpret_cast<const uint64_t*>(ctx0->hash_state) + 24); + } + // Optim - 99% time boundary - cn_explode_scratchpad<MEM, SOFT_AES, PREFETCH>((__m128i*)ctx0->hash_state, (__m128i*)ctx0->long_state); + cn_explode_scratchpad<MEM, SOFT_AES, PREFETCH, ALGO>((__m128i*)ctx0->hash_state, (__m128i*)ctx0->long_state); uint8_t* l0 = ctx0->long_state; uint64_t* h0 = (uint64_t*)ctx0->hash_state; @@ -315,8 +482,13 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c else cx = _mm_aesenc_si128(cx, _mm_set_epi64x(ah0, al0)); - _mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx)); + if(ALGO == cryptonight_monero) + cryptonight_monero_tweak((uint64_t*)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx)); + else + _mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx)); + idx0 = _mm_cvtsi128_si64(cx); + if(PREFETCH) _mm_prefetch((const char*)&l0[idx0 & MASK], _MM_HINT_T0); bx0 = cx; @@ -333,14 +505,28 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c if(PREFETCH) _mm_prefetch((const char*)&l0[al0 & MASK], _MM_HINT_T0); ah0 += lo; - ((uint64_t*)&l0[idx0 & MASK])[1] = ah0; + + if(ALGO == cryptonight_monero) + ((uint64_t*)&l0[idx0 & MASK])[1] = ah0 ^ monero_const; + else + ((uint64_t*)&l0[idx0 & MASK])[1] = ah0; ah0 ^= ch; idx0 = al0; + + if(ALGO == cryptonight_heavy) + { + int64_t n = ((int64_t*)&l0[idx0 & MASK])[0]; + int32_t d = ((int32_t*)&l0[idx0 & MASK])[2]; + int64_t q = n / (d | 0x5); + + ((int64_t*)&l0[idx0 & MASK])[0] = n ^ q; + idx0 = d ^ q; + } } // Optim - 90% time boundary - cn_implode_scratchpad<MEM, SOFT_AES, PREFETCH>((__m128i*)ctx0->long_state, (__m128i*)ctx0->hash_state); + cn_implode_scratchpad<MEM, SOFT_AES, PREFETCH, ALGO>((__m128i*)ctx0->long_state, (__m128i*)ctx0->hash_state); // Optim - 99% time boundary @@ -351,15 +537,34 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c // This lovely creation will do 2 cn hashes at a time. We have plenty of space on silicon // to fit temporary vars for two contexts. Function will read len*2 from input and write 64 bytes to output // We are still limited by L3 cache, so doubling will only work with CPUs where we have more than 2MB to core (Xeons) -template<size_t MASK, size_t ITERATIONS, size_t MEM, bool SOFT_AES, bool PREFETCH> +template<xmrstak_algo ALGO, bool SOFT_AES, bool PREFETCH> void cryptonight_double_hash(const void* input, size_t len, void* output, cryptonight_ctx** ctx) { + constexpr size_t MASK = cn_select_mask<ALGO>(); + constexpr size_t ITERATIONS = cn_select_iter<ALGO>(); + constexpr size_t MEM = cn_select_memory<ALGO>(); + + if(ALGO == cryptonight_monero && len < 43) + { + memset(output, 0, 64); + return; + } + keccak((const uint8_t *)input, len, ctx[0]->hash_state, 200); keccak((const uint8_t *)input+len, len, ctx[1]->hash_state, 200); + uint64_t monero_const_0, monero_const_1; + if(ALGO == cryptonight_monero) + { + monero_const_0 = *reinterpret_cast<const uint64_t*>(reinterpret_cast<const uint8_t*>(input) + 35); + monero_const_0 ^= *(reinterpret_cast<const uint64_t*>(ctx[0]->hash_state) + 24); + monero_const_1 = *reinterpret_cast<const uint64_t*>(reinterpret_cast<const uint8_t*>(input) + len + 35); + monero_const_1 ^= *(reinterpret_cast<const uint64_t*>(ctx[1]->hash_state) + 24); + } + // Optim - 99% time boundary - cn_explode_scratchpad<MEM, SOFT_AES, PREFETCH>((__m128i*)ctx[0]->hash_state, (__m128i*)ctx[0]->long_state); - cn_explode_scratchpad<MEM, SOFT_AES, PREFETCH>((__m128i*)ctx[1]->hash_state, (__m128i*)ctx[1]->long_state); + cn_explode_scratchpad<MEM, SOFT_AES, PREFETCH, ALGO>((__m128i*)ctx[0]->hash_state, (__m128i*)ctx[0]->long_state); + cn_explode_scratchpad<MEM, SOFT_AES, PREFETCH, ALGO>((__m128i*)ctx[1]->hash_state, (__m128i*)ctx[1]->long_state); uint8_t* l0 = ctx[0]->long_state; uint64_t* h0 = (uint64_t*)ctx[0]->hash_state; @@ -387,7 +592,11 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto else cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh0, axl0)); - _mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx)); + if(ALGO == cryptonight_monero) + cryptonight_monero_tweak((uint64_t*)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx)); + else + _mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx)); + idx0 = _mm_cvtsi128_si64(cx); bx0 = cx; @@ -401,7 +610,11 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto else cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh1, axl1)); - _mm_store_si128((__m128i *)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx)); + if(ALGO == cryptonight_monero) + cryptonight_monero_tweak((uint64_t*)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx)); + else + _mm_store_si128((__m128i *)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx)); + idx1 = _mm_cvtsi128_si64(cx); bx1 = cx; @@ -417,11 +630,26 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto axl0 += hi; axh0 += lo; ((uint64_t*)&l0[idx0 & MASK])[0] = axl0; - ((uint64_t*)&l0[idx0 & MASK])[1] = axh0; + + if(ALGO == cryptonight_monero) + ((uint64_t*)&l0[idx0 & MASK])[1] = axh0 ^ monero_const_0; + else + ((uint64_t*)&l0[idx0 & MASK])[1] = axh0; + axh0 ^= ch; axl0 ^= cl; idx0 = axl0; + if(ALGO == cryptonight_heavy) + { + int64_t n = ((int64_t*)&l0[idx0 & MASK])[0]; + int32_t d = ((int32_t*)&l0[idx0 & MASK])[2]; + int64_t q = n / (d | 0x5); + + ((int64_t*)&l0[idx0 & MASK])[0] = n ^ q; + idx0 = d ^ q; + } + if(PREFETCH) _mm_prefetch((const char*)&l0[idx0 & MASK], _MM_HINT_T0); @@ -433,18 +661,33 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto axl1 += hi; axh1 += lo; ((uint64_t*)&l1[idx1 & MASK])[0] = axl1; - ((uint64_t*)&l1[idx1 & MASK])[1] = axh1; + + if(ALGO == cryptonight_monero) + ((uint64_t*)&l1[idx1 & MASK])[1] = axh1 ^ monero_const_1; + else + ((uint64_t*)&l1[idx1 & MASK])[1] = axh1; + axh1 ^= ch; axl1 ^= cl; idx1 = axl1; + if(ALGO == cryptonight_heavy) + { + int64_t n = ((int64_t*)&l1[idx1 & MASK])[0]; + int32_t d = ((int32_t*)&l1[idx1 & MASK])[2]; + int64_t q = n / (d | 0x5); + + ((int64_t*)&l1[idx1 & MASK])[0] = n ^ q; + idx1 = d ^ q; + } + if(PREFETCH) _mm_prefetch((const char*)&l1[idx1 & MASK], _MM_HINT_T0); } // Optim - 90% time boundary - cn_implode_scratchpad<MEM, SOFT_AES, PREFETCH>((__m128i*)ctx[0]->long_state, (__m128i*)ctx[0]->hash_state); - cn_implode_scratchpad<MEM, SOFT_AES, PREFETCH>((__m128i*)ctx[1]->long_state, (__m128i*)ctx[1]->hash_state); + cn_implode_scratchpad<MEM, SOFT_AES, PREFETCH, ALGO>((__m128i*)ctx[0]->long_state, (__m128i*)ctx[0]->hash_state); + cn_implode_scratchpad<MEM, SOFT_AES, PREFETCH, ALGO>((__m128i*)ctx[1]->long_state, (__m128i*)ctx[1]->hash_state); // Optim - 99% time boundary @@ -455,12 +698,10 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto } #define CN_STEP1(a, b, c, l, ptr, idx) \ - a = _mm_xor_si128(a, c); \ - idx = _mm_cvtsi128_si64(a); \ ptr = (__m128i *)&l[idx & MASK]; \ if(PREFETCH) \ _mm_prefetch((const char*)ptr, _MM_HINT_T0); \ - c = _mm_load_si128(ptr) + c = _mm_load_si128(ptr); #define CN_STEP2(a, b, c, l, ptr, idx) \ if(SOFT_AES) \ @@ -468,30 +709,64 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto else \ c = _mm_aesenc_si128(c, a); \ b = _mm_xor_si128(b, c); \ - _mm_store_si128(ptr, b) + if(ALGO == cryptonight_monero) \ + cryptonight_monero_tweak((uint64_t*)ptr, b); \ + else \ + _mm_store_si128(ptr, b);\ #define CN_STEP3(a, b, c, l, ptr, idx) \ idx = _mm_cvtsi128_si64(c); \ ptr = (__m128i *)&l[idx & MASK]; \ if(PREFETCH) \ _mm_prefetch((const char*)ptr, _MM_HINT_T0); \ - b = _mm_load_si128(ptr) + b = _mm_load_si128(ptr); -#define CN_STEP4(a, b, c, l, ptr, idx) \ +#define CN_STEP4(a, b, c, l, mc, ptr, idx) \ lo = _umul128(idx, _mm_cvtsi128_si64(b), &hi); \ a = _mm_add_epi64(a, _mm_set_epi64x(lo, hi)); \ - _mm_store_si128(ptr, a) + if(ALGO == cryptonight_monero) \ + _mm_store_si128(ptr, _mm_xor_si128(a, mc)); \ + else \ + _mm_store_si128(ptr, a);\ + a = _mm_xor_si128(a, b); \ + idx = _mm_cvtsi128_si64(a); \ + if(ALGO == cryptonight_heavy) \ + { \ + int64_t n = ((int64_t*)&l[idx & MASK])[0]; \ + int32_t d = ((int32_t*)&l[idx & MASK])[2]; \ + int64_t q = n / (d | 0x5); \ + ((int64_t*)&l[idx & MASK])[0] = n ^ q; \ + idx = d ^ q; \ + } + +#define CONST_INIT(ctx, n) \ + __m128i mc##n = _mm_set_epi64x(*reinterpret_cast<const uint64_t*>(reinterpret_cast<const uint8_t*>(input) + n * len + 35) ^ \ + *(reinterpret_cast<const uint64_t*>((ctx)->hash_state) + 24), 0); // This lovelier creation will do 3 cn hashes at a time. -template<size_t MASK, size_t ITERATIONS, size_t MEM, bool SOFT_AES, bool PREFETCH> +template<xmrstak_algo ALGO, bool SOFT_AES, bool PREFETCH> void cryptonight_triple_hash(const void* input, size_t len, void* output, cryptonight_ctx** ctx) { + constexpr size_t MASK = cn_select_mask<ALGO>(); + constexpr size_t ITERATIONS = cn_select_iter<ALGO>(); + constexpr size_t MEM = cn_select_memory<ALGO>(); + + if(ALGO == cryptonight_monero && len < 43) + { + memset(output, 0, 32 * 3); + return; + } + for (size_t i = 0; i < 3; i++) { keccak((const uint8_t *)input + len * i, len, ctx[i]->hash_state, 200); - cn_explode_scratchpad<MEM, SOFT_AES, PREFETCH>((__m128i*)ctx[i]->hash_state, (__m128i*)ctx[i]->long_state); + cn_explode_scratchpad<MEM, SOFT_AES, PREFETCH, ALGO>((__m128i*)ctx[i]->hash_state, (__m128i*)ctx[i]->long_state); } + CONST_INIT(ctx[0], 0); + CONST_INIT(ctx[1], 1); + CONST_INIT(ctx[2], 2); + uint8_t* l0 = ctx[0]->long_state; uint64_t* h0 = (uint64_t*)ctx[0]->hash_state; uint8_t* l1 = ctx[1]->long_state; @@ -509,9 +784,14 @@ void cryptonight_triple_hash(const void* input, size_t len, void* output, crypto __m128i cx1 = _mm_set_epi64x(0, 0); __m128i cx2 = _mm_set_epi64x(0, 0); + uint64_t idx0, idx1, idx2; + idx0 = _mm_cvtsi128_si64(ax0); + idx1 = _mm_cvtsi128_si64(ax1); + idx2 = _mm_cvtsi128_si64(ax2); + for (size_t i = 0; i < ITERATIONS/2; i++) { - uint64_t idx0, idx1, idx2, hi, lo; + uint64_t hi, lo; __m128i *ptr0, *ptr1, *ptr2; // EVEN ROUND @@ -527,9 +807,9 @@ void cryptonight_triple_hash(const void* input, size_t len, void* output, crypto CN_STEP3(ax1, bx1, cx1, l1, ptr1, idx1); CN_STEP3(ax2, bx2, cx2, l2, ptr2, idx2); - CN_STEP4(ax0, bx0, cx0, l0, ptr0, idx0); - CN_STEP4(ax1, bx1, cx1, l1, ptr1, idx1); - CN_STEP4(ax2, bx2, cx2, l2, ptr2, idx2); + CN_STEP4(ax0, bx0, cx0, l0, mc0, ptr0, idx0); + CN_STEP4(ax1, bx1, cx1, l1, mc1, ptr1, idx1); + CN_STEP4(ax2, bx2, cx2, l2, mc2, ptr2, idx2); // ODD ROUND CN_STEP1(ax0, cx0, bx0, l0, ptr0, idx0); @@ -544,29 +824,44 @@ void cryptonight_triple_hash(const void* input, size_t len, void* output, crypto CN_STEP3(ax1, cx1, bx1, l1, ptr1, idx1); CN_STEP3(ax2, cx2, bx2, l2, ptr2, idx2); - CN_STEP4(ax0, cx0, bx0, l0, ptr0, idx0); - CN_STEP4(ax1, cx1, bx1, l1, ptr1, idx1); - CN_STEP4(ax2, cx2, bx2, l2, ptr2, idx2); + CN_STEP4(ax0, cx0, bx0, l0, mc0, ptr0, idx0); + CN_STEP4(ax1, cx1, bx1, l1, mc1, ptr1, idx1); + CN_STEP4(ax2, cx2, bx2, l2, mc2, ptr2, idx2); } for (size_t i = 0; i < 3; i++) { - cn_implode_scratchpad<MEM, SOFT_AES, PREFETCH>((__m128i*)ctx[i]->long_state, (__m128i*)ctx[i]->hash_state); + cn_implode_scratchpad<MEM, SOFT_AES, PREFETCH, ALGO>((__m128i*)ctx[i]->long_state, (__m128i*)ctx[i]->hash_state); keccakf((uint64_t*)ctx[i]->hash_state, 24); extra_hashes[ctx[i]->hash_state[0] & 3](ctx[i]->hash_state, 200, (char*)output + 32 * i); } } // This even lovelier creation will do 4 cn hashes at a time. -template<size_t MASK, size_t ITERATIONS, size_t MEM, bool SOFT_AES, bool PREFETCH> +template<xmrstak_algo ALGO, bool SOFT_AES, bool PREFETCH> void cryptonight_quad_hash(const void* input, size_t len, void* output, cryptonight_ctx** ctx) { + constexpr size_t MASK = cn_select_mask<ALGO>(); + constexpr size_t ITERATIONS = cn_select_iter<ALGO>(); + constexpr size_t MEM = cn_select_memory<ALGO>(); + + if(ALGO == cryptonight_monero && len < 43) + { + memset(output, 0, 32 * 4); + return; + } + for (size_t i = 0; i < 4; i++) { keccak((const uint8_t *)input + len * i, len, ctx[i]->hash_state, 200); - cn_explode_scratchpad<MEM, SOFT_AES, PREFETCH>((__m128i*)ctx[i]->hash_state, (__m128i*)ctx[i]->long_state); + cn_explode_scratchpad<MEM, SOFT_AES, PREFETCH, ALGO>((__m128i*)ctx[i]->hash_state, (__m128i*)ctx[i]->long_state); } + CONST_INIT(ctx[0], 0); + CONST_INIT(ctx[1], 1); + CONST_INIT(ctx[2], 2); + CONST_INIT(ctx[3], 3); + uint8_t* l0 = ctx[0]->long_state; uint64_t* h0 = (uint64_t*)ctx[0]->hash_state; uint8_t* l1 = ctx[1]->long_state; @@ -588,10 +883,16 @@ void cryptonight_quad_hash(const void* input, size_t len, void* output, cryptoni __m128i cx1 = _mm_set_epi64x(0, 0); __m128i cx2 = _mm_set_epi64x(0, 0); __m128i cx3 = _mm_set_epi64x(0, 0); - + + uint64_t idx0, idx1, idx2, idx3; + idx0 = _mm_cvtsi128_si64(ax0); + idx1 = _mm_cvtsi128_si64(ax1); + idx2 = _mm_cvtsi128_si64(ax2); + idx3 = _mm_cvtsi128_si64(ax3); + for (size_t i = 0; i < ITERATIONS/2; i++) { - uint64_t idx0, idx1, idx2, idx3, hi, lo; + uint64_t hi, lo; __m128i *ptr0, *ptr1, *ptr2, *ptr3; // EVEN ROUND @@ -610,10 +911,10 @@ void cryptonight_quad_hash(const void* input, size_t len, void* output, cryptoni CN_STEP3(ax2, bx2, cx2, l2, ptr2, idx2); CN_STEP3(ax3, bx3, cx3, l3, ptr3, idx3); - CN_STEP4(ax0, bx0, cx0, l0, ptr0, idx0); - CN_STEP4(ax1, bx1, cx1, l1, ptr1, idx1); - CN_STEP4(ax2, bx2, cx2, l2, ptr2, idx2); - CN_STEP4(ax3, bx3, cx3, l3, ptr3, idx3); + CN_STEP4(ax0, bx0, cx0, l0, mc0, ptr0, idx0); + CN_STEP4(ax1, bx1, cx1, l1, mc1, ptr1, idx1); + CN_STEP4(ax2, bx2, cx2, l2, mc2, ptr2, idx2); + CN_STEP4(ax3, bx3, cx3, l3, mc3, ptr3, idx3); // ODD ROUND CN_STEP1(ax0, cx0, bx0, l0, ptr0, idx0); @@ -631,30 +932,46 @@ void cryptonight_quad_hash(const void* input, size_t len, void* output, cryptoni CN_STEP3(ax2, cx2, bx2, l2, ptr2, idx2); CN_STEP3(ax3, cx3, bx3, l3, ptr3, idx3); - CN_STEP4(ax0, cx0, bx0, l0, ptr0, idx0); - CN_STEP4(ax1, cx1, bx1, l1, ptr1, idx1); - CN_STEP4(ax2, cx2, bx2, l2, ptr2, idx2); - CN_STEP4(ax3, cx3, bx3, l3, ptr3, idx3); + CN_STEP4(ax0, cx0, bx0, l0, mc0, ptr0, idx0); + CN_STEP4(ax1, cx1, bx1, l1, mc1, ptr1, idx1); + CN_STEP4(ax2, cx2, bx2, l2, mc2, ptr2, idx2); + CN_STEP4(ax3, cx3, bx3, l3, mc3, ptr3, idx3); } for (size_t i = 0; i < 4; i++) { - cn_implode_scratchpad<MEM, SOFT_AES, PREFETCH>((__m128i*)ctx[i]->long_state, (__m128i*)ctx[i]->hash_state); + cn_implode_scratchpad<MEM, SOFT_AES, PREFETCH, ALGO>((__m128i*)ctx[i]->long_state, (__m128i*)ctx[i]->hash_state); keccakf((uint64_t*)ctx[i]->hash_state, 24); extra_hashes[ctx[i]->hash_state[0] & 3](ctx[i]->hash_state, 200, (char*)output + 32 * i); } } // This most lovely creation will do 5 cn hashes at a time. -template<size_t MASK, size_t ITERATIONS, size_t MEM, bool SOFT_AES, bool PREFETCH> +template<xmrstak_algo ALGO, bool SOFT_AES, bool PREFETCH> void cryptonight_penta_hash(const void* input, size_t len, void* output, cryptonight_ctx** ctx) { + constexpr size_t MASK = cn_select_mask<ALGO>(); + constexpr size_t ITERATIONS = cn_select_iter<ALGO>(); + constexpr size_t MEM = cn_select_memory<ALGO>(); + + if(ALGO == cryptonight_monero && len < 43) + { + memset(output, 0, 32 * 5); + return; + } + for (size_t i = 0; i < 5; i++) { keccak((const uint8_t *)input + len * i, len, ctx[i]->hash_state, 200); - cn_explode_scratchpad<MEM, SOFT_AES, PREFETCH>((__m128i*)ctx[i]->hash_state, (__m128i*)ctx[i]->long_state); + cn_explode_scratchpad<MEM, SOFT_AES, PREFETCH, ALGO>((__m128i*)ctx[i]->hash_state, (__m128i*)ctx[i]->long_state); } + CONST_INIT(ctx[0], 0); + CONST_INIT(ctx[1], 1); + CONST_INIT(ctx[2], 2); + CONST_INIT(ctx[3], 3); + CONST_INIT(ctx[4], 4); + uint8_t* l0 = ctx[0]->long_state; uint64_t* h0 = (uint64_t*)ctx[0]->hash_state; uint8_t* l1 = ctx[1]->long_state; @@ -682,9 +999,16 @@ void cryptonight_penta_hash(const void* input, size_t len, void* output, crypton __m128i cx3 = _mm_set_epi64x(0, 0); __m128i cx4 = _mm_set_epi64x(0, 0); + uint64_t idx0, idx1, idx2, idx3, idx4; + idx0 = _mm_cvtsi128_si64(ax0); + idx1 = _mm_cvtsi128_si64(ax1); + idx2 = _mm_cvtsi128_si64(ax2); + idx3 = _mm_cvtsi128_si64(ax3); + idx4 = _mm_cvtsi128_si64(ax4); + for (size_t i = 0; i < ITERATIONS/2; i++) { - uint64_t idx0, idx1, idx2, idx3, idx4, hi, lo; + uint64_t hi, lo; __m128i *ptr0, *ptr1, *ptr2, *ptr3, *ptr4; // EVEN ROUND @@ -706,11 +1030,11 @@ void cryptonight_penta_hash(const void* input, size_t len, void* output, crypton CN_STEP3(ax3, bx3, cx3, l3, ptr3, idx3); CN_STEP3(ax4, bx4, cx4, l4, ptr4, idx4); - CN_STEP4(ax0, bx0, cx0, l0, ptr0, idx0); - CN_STEP4(ax1, bx1, cx1, l1, ptr1, idx1); - CN_STEP4(ax2, bx2, cx2, l2, ptr2, idx2); - CN_STEP4(ax3, bx3, cx3, l3, ptr3, idx3); - CN_STEP4(ax4, bx4, cx4, l4, ptr4, idx4); + CN_STEP4(ax0, bx0, cx0, l0, mc0, ptr0, idx0); + CN_STEP4(ax1, bx1, cx1, l1, mc1, ptr1, idx1); + CN_STEP4(ax2, bx2, cx2, l2, mc2, ptr2, idx2); + CN_STEP4(ax3, bx3, cx3, l3, mc3, ptr3, idx3); + CN_STEP4(ax4, bx4, cx4, l4, mc4, ptr4, idx4); // ODD ROUND CN_STEP1(ax0, cx0, bx0, l0, ptr0, idx0); @@ -731,16 +1055,16 @@ void cryptonight_penta_hash(const void* input, size_t len, void* output, crypton CN_STEP3(ax3, cx3, bx3, l3, ptr3, idx3); CN_STEP3(ax4, cx4, bx4, l4, ptr4, idx4); - CN_STEP4(ax0, cx0, bx0, l0, ptr0, idx0); - CN_STEP4(ax1, cx1, bx1, l1, ptr1, idx1); - CN_STEP4(ax2, cx2, bx2, l2, ptr2, idx2); - CN_STEP4(ax3, cx3, bx3, l3, ptr3, idx3); - CN_STEP4(ax4, cx4, bx4, l4, ptr4, idx4); + CN_STEP4(ax0, cx0, bx0, l0, mc0, ptr0, idx0); + CN_STEP4(ax1, cx1, bx1, l1, mc1, ptr1, idx1); + CN_STEP4(ax2, cx2, bx2, l2, mc2, ptr2, idx2); + CN_STEP4(ax3, cx3, bx3, l3, mc3, ptr3, idx3); + CN_STEP4(ax4, cx4, bx4, l4, mc4, ptr4, idx4); } for (size_t i = 0; i < 5; i++) { - cn_implode_scratchpad<MEM, SOFT_AES, PREFETCH>((__m128i*)ctx[i]->long_state, (__m128i*)ctx[i]->hash_state); + cn_implode_scratchpad<MEM, SOFT_AES, PREFETCH, ALGO>((__m128i*)ctx[i]->long_state, (__m128i*)ctx[i]->hash_state); keccakf((uint64_t*)ctx[i]->hash_state, 24); extra_hashes[ctx[i]->hash_state[0] & 3](ctx[i]->hash_state, 200, (char*)output + 32 * i); } diff --git a/xmrstak/backend/cpu/crypto/cryptonight_common.cpp b/xmrstak/backend/cpu/crypto/cryptonight_common.cpp index 1026b04..17fa24b 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight_common.cpp +++ b/xmrstak/backend/cpu/crypto/cryptonight_common.cpp @@ -28,9 +28,9 @@ extern "C" #include "c_jh.h" #include "c_skein.h" } +#include "xmrstak/backend/cryptonight.hpp" #include "cryptonight.h" #include "cryptonight_aesni.h" -#include "xmrstak/backend/cryptonight.hpp" #include "xmrstak/misc/console.hpp" #include "xmrstak/jconf.hpp" #include <stdio.h> @@ -202,15 +202,8 @@ size_t cryptonight_init(size_t use_fast_mem, size_t use_mlock, alloc_msg* msg) cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, alloc_msg* msg) { - size_t hashMemSize; - if(::jconf::inst()->IsCurrencyMonero()) - { - hashMemSize = MONERO_MEMORY; - } - else - { - hashMemSize = AEON_MEMORY; - } + size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo()); + cryptonight_ctx* ptr = (cryptonight_ctx*)_mm_malloc(sizeof(cryptonight_ctx), 4096); if(use_fast_mem == 0) @@ -285,15 +278,8 @@ cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, al void cryptonight_free_ctx(cryptonight_ctx* ctx) { - size_t hashMemSize; - if(::jconf::inst()->IsCurrencyMonero()) - { - hashMemSize = MONERO_MEMORY; - } - else - { - hashMemSize = AEON_MEMORY; - } + size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo()); + if(ctx->ctx_info[0] != 0) { #ifdef _WIN32 diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp index cef4f8e..e263aca 100644 --- a/xmrstak/backend/cpu/minethd.cpp +++ b/xmrstak/backend/cpu/minethd.cpp @@ -231,45 +231,44 @@ bool minethd::self_test() bool bResult = true; - bool mineMonero = ::jconf::inst()->IsCurrencyMonero(); - if(mineMonero) + if(::jconf::inst()->GetMiningAlgo() == cryptonight) { unsigned char out[32 * MAX_N]; cn_hash_fun hashf; cn_hash_fun_multi hashf_multi; - hashf = func_selector(::jconf::inst()->HaveHardwareAes(), false, mineMonero); + hashf = func_selector(::jconf::inst()->HaveHardwareAes(), false, xmrstak_algo::cryptonight); hashf("This is a test", 14, out, ctx[0]); bResult = memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 32) == 0; - hashf = func_selector(::jconf::inst()->HaveHardwareAes(), true, mineMonero); + hashf = func_selector(::jconf::inst()->HaveHardwareAes(), true, xmrstak_algo::cryptonight); hashf("This is a test", 14, out, ctx[0]); bResult &= memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 32) == 0; - hashf_multi = func_multi_selector(2, ::jconf::inst()->HaveHardwareAes(), false, mineMonero); + hashf_multi = func_multi_selector(2, ::jconf::inst()->HaveHardwareAes(), false, xmrstak_algo::cryptonight); hashf_multi("The quick brown fox jumps over the lazy dogThe quick brown fox jumps over the lazy log", 43, out, ctx); bResult &= memcmp(out, "\x3e\xbb\x7f\x9f\x7d\x27\x3d\x7c\x31\x8d\x86\x94\x77\x55\x0c\xc8\x00\xcf\xb1\x1b\x0c\xad\xb7\xff\xbd\xf6\xf8\x9f\x3a\x47\x1c\x59" "\xb4\x77\xd5\x02\xe4\xd8\x48\x7f\x42\xdf\xe3\x8e\xed\x73\x81\x7a\xda\x91\xb7\xe2\x63\xd2\x91\x71\xb6\x5c\x44\x3a\x01\x2a\x41\x22", 64) == 0; - hashf_multi = func_multi_selector(2, ::jconf::inst()->HaveHardwareAes(), true, mineMonero); + hashf_multi = func_multi_selector(2, ::jconf::inst()->HaveHardwareAes(), true, xmrstak_algo::cryptonight); hashf_multi("The quick brown fox jumps over the lazy dogThe quick brown fox jumps over the lazy log", 43, out, ctx); bResult &= memcmp(out, "\x3e\xbb\x7f\x9f\x7d\x27\x3d\x7c\x31\x8d\x86\x94\x77\x55\x0c\xc8\x00\xcf\xb1\x1b\x0c\xad\xb7\xff\xbd\xf6\xf8\x9f\x3a\x47\x1c\x59" "\xb4\x77\xd5\x02\xe4\xd8\x48\x7f\x42\xdf\xe3\x8e\xed\x73\x81\x7a\xda\x91\xb7\xe2\x63\xd2\x91\x71\xb6\x5c\x44\x3a\x01\x2a\x41\x22", 64) == 0; - hashf_multi = func_multi_selector(3, ::jconf::inst()->HaveHardwareAes(), false, mineMonero); + hashf_multi = func_multi_selector(3, ::jconf::inst()->HaveHardwareAes(), false, xmrstak_algo::cryptonight); hashf_multi("This is a testThis is a testThis is a test", 14, out, ctx); bResult &= memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05" "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05" "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 96) == 0; - hashf_multi = func_multi_selector(4, ::jconf::inst()->HaveHardwareAes(), false, mineMonero); + hashf_multi = func_multi_selector(4, ::jconf::inst()->HaveHardwareAes(), false, xmrstak_algo::cryptonight); hashf_multi("This is a testThis is a testThis is a testThis is a test", 14, out, ctx); bResult &= memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05" "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05" "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05" "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 128) == 0; - hashf_multi = func_multi_selector(5, ::jconf::inst()->HaveHardwareAes(), false, mineMonero); + hashf_multi = func_multi_selector(5, ::jconf::inst()->HaveHardwareAes(), false, xmrstak_algo::cryptonight); hashf_multi("This is a testThis is a testThis is a testThis is a testThis is a test", 14, out, ctx); bResult &= memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05" "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05" @@ -277,6 +276,12 @@ bool minethd::self_test() "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05" "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 160) == 0; } + else if(::jconf::inst()->GetMiningAlgo() == cryptonight_lite) + { + } + else if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero) + { + } for (int i = 0; i < MAX_N; i++) cryptonight_free_ctx(ctx[i]); @@ -340,48 +345,56 @@ void minethd::consume_work() globalStates::inst().inst().iConsumeCnt++; } -minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, bool mineMonero) +minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmrstak_algo algo) { // We have two independent flag bits in the functions // therefore we will build a binary digit and select the // function as a two digit binary - // Digit order SOFT_AES, NO_PREFETCH, MINER_ALGO + + uint8_t algv; + switch(algo) + { + case cryptonight: + algv = 2; + break; + case cryptonight_lite: + algv = 1; + break; + case cryptonight_monero: + algv = 0; + break; + case cryptonight_heavy: + algv = 3; + break; + default: + algv = 2; + break; + } static const cn_hash_fun func_table[] = { - /* there will be 8 function entries if `CONF_NO_MONERO` and `CONF_NO_AEON` - * is not defined. If one is defined there will be 4 entries. - */ -#ifndef CONF_NO_MONERO - cryptonight_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, false>, - cryptonight_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, true>, - cryptonight_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, false>, - cryptonight_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, true> -#endif -#if (!defined(CONF_NO_AEON)) && (!defined(CONF_NO_MONERO)) - // comma will be added only if Monero and Aeon is build - , -#endif -#ifndef CONF_NO_AEON - cryptonight_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, false>, - cryptonight_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, true>, - cryptonight_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, false>, - cryptonight_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, true> -#endif + cryptonight_hash<cryptonight_monero, false, false>, + cryptonight_hash<cryptonight_monero, true, false>, + cryptonight_hash<cryptonight_monero, false, true>, + cryptonight_hash<cryptonight_monero, true, true>, + cryptonight_hash<cryptonight_lite, false, false>, + cryptonight_hash<cryptonight_lite, true, false>, + cryptonight_hash<cryptonight_lite, false, true>, + cryptonight_hash<cryptonight_lite, true, true>, + cryptonight_hash<cryptonight, false, false>, + cryptonight_hash<cryptonight, true, false>, + cryptonight_hash<cryptonight, false, true>, + cryptonight_hash<cryptonight, true, true>, + cryptonight_hash<cryptonight_heavy, false, false>, + cryptonight_hash<cryptonight_heavy, true, false>, + cryptonight_hash<cryptonight_heavy, false, true>, + cryptonight_hash<cryptonight_heavy, true, true> }; - std::bitset<3> digit; - digit.set(0, !bNoPrefetch); - digit.set(1, !bHaveAes); - - // define aeon settings -#if defined(CONF_NO_AEON) || defined(CONF_NO_MONERO) - // ignore 3rd bit if only one currency is active - digit.set(2, 0); -#else - digit.set(2, !mineMonero); -#endif + std::bitset<2> digit; + digit.set(0, !bHaveAes); + digit.set(1, !bNoPrefetch); - return func_table[digit.to_ulong()]; + return func_table[ algv << 2 | digit.to_ulong() ]; } void minethd::work_main() @@ -401,7 +414,7 @@ void minethd::work_main() uint32_t* piNonce; job_result result; - hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->IsCurrencyMonero()); + hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo()); ctx = minethd_alloc_ctx(); piHashVal = (uint64_t*)(result.bResult + 24); @@ -434,6 +447,22 @@ void minethd::work_main() if(oWork.bNiceHash) result.iNonce = *piNonce; + if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero) + { + if(oWork.bWorkBlob[0] >= 7) + hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight_monero); + else + hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight); + } + + if(::jconf::inst()->GetMiningAlgo() == cryptonight_heavy) + { + if(oWork.bWorkBlob[0] >= 3) + hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight_heavy); + else + hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight); + } + while(globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo) { if ((iCount++ & 0xF) == 0) //Store stats every 16 hashes @@ -465,93 +494,105 @@ void minethd::work_main() cryptonight_free_ctx(ctx); } -minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, bool bNoPrefetch, bool mineMonero) +minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, bool bNoPrefetch, xmrstak_algo algo) { // We have two independent flag bits in the functions // therefore we will build a binary digit and select the // function as a two digit binary - // Digit order SOFT_AES, NO_PREFETCH + + uint8_t algv; + switch(algo) + { + case cryptonight: + algv = 2; + break; + case cryptonight_lite: + algv = 1; + break; + case cryptonight_monero: + algv = 0; + break; + default: + algv = 2; + break; + } static const cn_hash_fun_multi func_table[] = { - /* there will be 8*(MAX_N-1) function entries if `CONF_NO_MONERO` and `CONF_NO_AEON` - * is not defined. If one is defined there will be 4*(MAX_N-1) entries. - */ -#ifndef CONF_NO_MONERO - cryptonight_double_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, false>, - cryptonight_double_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, true>, - cryptonight_double_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, false>, - cryptonight_double_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, true>, - cryptonight_triple_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, false>, - cryptonight_triple_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, true>, - cryptonight_triple_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, false>, - cryptonight_triple_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, true>, - cryptonight_quad_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, false>, - cryptonight_quad_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, true>, - cryptonight_quad_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, false>, - cryptonight_quad_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, true>, - cryptonight_penta_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, false>, - cryptonight_penta_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, false, true>, - cryptonight_penta_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, false>, - cryptonight_penta_hash<MONERO_MASK, MONERO_ITER, MONERO_MEMORY, true, true> -#endif -#if (!defined(CONF_NO_AEON)) && (!defined(CONF_NO_MONERO)) - // comma will be added only if Monero and Aeon is build - , -#endif -#ifndef CONF_NO_AEON - cryptonight_double_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, false>, - cryptonight_double_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, true>, - cryptonight_double_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, false>, - cryptonight_double_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, true>, - cryptonight_triple_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, false>, - cryptonight_triple_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, true>, - cryptonight_triple_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, false>, - cryptonight_triple_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, true>, - cryptonight_quad_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, false>, - cryptonight_quad_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, true>, - cryptonight_quad_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, false>, - cryptonight_quad_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, true>, - cryptonight_penta_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, false>, - cryptonight_penta_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, false, true>, - cryptonight_penta_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, false>, - cryptonight_penta_hash<AEON_MASK, AEON_ITER, AEON_MEMORY, true, true> -#endif + cryptonight_double_hash<cryptonight_monero, false, false>, + cryptonight_double_hash<cryptonight_monero, true, false>, + cryptonight_double_hash<cryptonight_monero, false, true>, + cryptonight_double_hash<cryptonight_monero, true, true>, + cryptonight_triple_hash<cryptonight_monero, false, false>, + cryptonight_triple_hash<cryptonight_monero, true, false>, + cryptonight_triple_hash<cryptonight_monero, false, true>, + cryptonight_triple_hash<cryptonight_monero, true, true>, + cryptonight_quad_hash<cryptonight_monero, false, false>, + cryptonight_quad_hash<cryptonight_monero, true, false>, + cryptonight_quad_hash<cryptonight_monero, false, true>, + cryptonight_quad_hash<cryptonight_monero, true, true>, + cryptonight_penta_hash<cryptonight_monero, false, false>, + cryptonight_penta_hash<cryptonight_monero, true, false>, + cryptonight_penta_hash<cryptonight_monero, false, true>, + cryptonight_penta_hash<cryptonight_monero, true, true>, + cryptonight_double_hash<cryptonight_lite, false, false>, + cryptonight_double_hash<cryptonight_lite, true, false>, + cryptonight_double_hash<cryptonight_lite, false, true>, + cryptonight_double_hash<cryptonight_lite, true, true>, + cryptonight_triple_hash<cryptonight_lite, false, false>, + cryptonight_triple_hash<cryptonight_lite, true, false>, + cryptonight_triple_hash<cryptonight_lite, false, true>, + cryptonight_triple_hash<cryptonight_lite, true, true>, + cryptonight_quad_hash<cryptonight_lite, false, false>, + cryptonight_quad_hash<cryptonight_lite, true, false>, + cryptonight_quad_hash<cryptonight_lite, false, true>, + cryptonight_quad_hash<cryptonight_lite, true, true>, + cryptonight_penta_hash<cryptonight_lite, false, false>, + cryptonight_penta_hash<cryptonight_lite, true, false>, + cryptonight_penta_hash<cryptonight_lite, false, true>, + cryptonight_penta_hash<cryptonight_lite, true, true>, + cryptonight_double_hash<cryptonight, false, false>, + cryptonight_double_hash<cryptonight, true, false>, + cryptonight_double_hash<cryptonight, false, true>, + cryptonight_double_hash<cryptonight, true, true>, + cryptonight_triple_hash<cryptonight, false, false>, + cryptonight_triple_hash<cryptonight, true, false>, + cryptonight_triple_hash<cryptonight, false, true>, + cryptonight_triple_hash<cryptonight, true, true>, + cryptonight_quad_hash<cryptonight, false, false>, + cryptonight_quad_hash<cryptonight, true, false>, + cryptonight_quad_hash<cryptonight, false, true>, + cryptonight_quad_hash<cryptonight, true, true>, + cryptonight_penta_hash<cryptonight, false, false>, + cryptonight_penta_hash<cryptonight, true, false>, + cryptonight_penta_hash<cryptonight, false, true>, + cryptonight_penta_hash<cryptonight, true, true> }; std::bitset<2> digit; - digit.set(0, !bNoPrefetch); - digit.set(1, !bHaveAes); - - // define aeon settings -#if defined(CONF_NO_AEON) || defined(CONF_NO_MONERO) - // ignore miner algo if only one currency is active - size_t miner_algo_base = 0; -#else - size_t miner_algo_base = mineMonero ? 0 : 4*(MAX_N-1); -#endif - - N = (N<2) ? 2 : (N>MAX_N) ? MAX_N : N; - return func_table[miner_algo_base + 4*(N-2) + digit.to_ulong()]; + digit.set(0, !bHaveAes); + digit.set(1, !bNoPrefetch); + + return func_table[algv << 4 | (N-2) << 2 | digit.to_ulong()]; } void minethd::double_work_main() { - multiway_work_main<2>(func_multi_selector(2, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->IsCurrencyMonero())); + multiway_work_main<2>(func_multi_selector(2, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo())); } void minethd::triple_work_main() { - multiway_work_main<3>(func_multi_selector(3, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->IsCurrencyMonero())); + multiway_work_main<3>(func_multi_selector(3, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo())); } void minethd::quad_work_main() { - multiway_work_main<4>(func_multi_selector(4, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->IsCurrencyMonero())); + multiway_work_main<4>(func_multi_selector(4, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo())); } void minethd::penta_work_main() { - multiway_work_main<5>(func_multi_selector(5, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->IsCurrencyMonero())); + multiway_work_main<5>(func_multi_selector(5, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo())); } template<size_t N> @@ -621,6 +662,22 @@ void minethd::multiway_work_main(cn_hash_fun_multi hash_fun_multi) if(oWork.bNiceHash) iNonce = *piNonce[0]; + if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero) + { + if(oWork.bWorkBlob[0] >= 7) + hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight_monero); + else + hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight); + } + + if(::jconf::inst()->GetMiningAlgo() == cryptonight_heavy) + { + if(oWork.bWorkBlob[0] >= 3) + hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight_heavy); + else + hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, cryptonight); + } + while (globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo) { if ((iCount++ & 0x7) == 0) //Store stats every 8*N hashes diff --git a/xmrstak/backend/cpu/minethd.hpp b/xmrstak/backend/cpu/minethd.hpp index 0433d0d..ef1bbd2 100644 --- a/xmrstak/backend/cpu/minethd.hpp +++ b/xmrstak/backend/cpu/minethd.hpp @@ -1,5 +1,6 @@ #pragma once +#include "xmrstak/jconf.hpp" #include "crypto/cryptonight.h" #include "xmrstak/backend/miner_work.hpp" #include "xmrstak/backend/iBackend.hpp" @@ -23,14 +24,14 @@ public: typedef void (*cn_hash_fun)(const void*, size_t, void*, cryptonight_ctx*); - static cn_hash_fun func_selector(bool bHaveAes, bool bNoPrefetch, bool mineMonero); + static cn_hash_fun func_selector(bool bHaveAes, bool bNoPrefetch, xmrstak_algo algo); static bool thd_setaffinity(std::thread::native_handle_type h, uint64_t cpu_id); static cryptonight_ctx* minethd_alloc_ctx(); private: typedef void (*cn_hash_fun_multi)(const void*, size_t, void*, cryptonight_ctx**); - static cn_hash_fun_multi func_multi_selector(size_t N, bool bHaveAes, bool bNoPrefetch, bool mineMonero); + static cn_hash_fun_multi func_multi_selector(size_t N, bool bHaveAes, bool bNoPrefetch, xmrstak_algo algo); minethd(miner_work& pWork, size_t iNo, int iMultiway, bool no_prefetch, int64_t affinity); diff --git a/xmrstak/backend/cryptonight.hpp b/xmrstak/backend/cryptonight.hpp index 0ef5ae7..fe10a9f 100644 --- a/xmrstak/backend/cryptonight.hpp +++ b/xmrstak/backend/cryptonight.hpp @@ -1,12 +1,123 @@ #pragma once +#include <stddef.h> +#include <inttypes.h> +#include <type_traits> + +enum xmrstak_algo +{ + invalid_algo = 0, + cryptonight = 1, + cryptonight_lite = 2, + cryptonight_monero = 3, + cryptonight_heavy = 4 +}; // define aeon settings -#define AEON_MEMORY 1048576llu -#define AEON_MASK 0xFFFF0 -#define AEON_ITER 0x40000 +constexpr size_t CRYPTONIGHT_LITE_MEMORY = 1 * 1024 * 1024; +constexpr uint32_t CRYPTONIGHT_LITE_MASK = 0xFFFF0; +constexpr uint32_t CRYPTONIGHT_LITE_ITER = 0x40000; + +constexpr size_t CRYPTONIGHT_MEMORY = 2 * 1024 * 1024; +constexpr uint32_t CRYPTONIGHT_MASK = 0x1FFFF0; +constexpr uint32_t CRYPTONIGHT_ITER = 0x80000; + +constexpr size_t CRYPTONIGHT_HEAVY_MEMORY = 4 * 1024 * 1024; +constexpr uint32_t CRYPTONIGHT_HEAVY_MASK = 0x3FFFF0; +constexpr uint32_t CRYPTONIGHT_HEAVY_ITER = 0x40000; + +template<xmrstak_algo ALGO> +inline constexpr size_t cn_select_memory() { return 0; } + +template<> +inline constexpr size_t cn_select_memory<cryptonight>() { return CRYPTONIGHT_MEMORY; } + +template<> +inline constexpr size_t cn_select_memory<cryptonight_lite>() { return CRYPTONIGHT_LITE_MEMORY; } + +template<> +inline constexpr size_t cn_select_memory<cryptonight_monero>() { return CRYPTONIGHT_MEMORY; } + +template<> +inline constexpr size_t cn_select_memory<cryptonight_heavy>() { return CRYPTONIGHT_HEAVY_MEMORY; } + + +inline size_t cn_select_memory(xmrstak_algo algo) +{ + switch(algo) + { + case cryptonight: + return CRYPTONIGHT_MEMORY; + case cryptonight_lite: + return CRYPTONIGHT_LITE_MEMORY; + case cryptonight_monero: + return CRYPTONIGHT_MEMORY; + case cryptonight_heavy: + return CRYPTONIGHT_HEAVY_MEMORY; + default: + return 0; + } +} + +template<xmrstak_algo ALGO> +inline constexpr uint32_t cn_select_mask() { return 0; } + +template<> +inline constexpr uint32_t cn_select_mask<cryptonight>() { return CRYPTONIGHT_MASK; } + +template<> +inline constexpr uint32_t cn_select_mask<cryptonight_lite>() { return CRYPTONIGHT_LITE_MASK; } + +template<> +inline constexpr uint32_t cn_select_mask<cryptonight_monero>() { return CRYPTONIGHT_MASK; } + +template<> +inline constexpr uint32_t cn_select_mask<cryptonight_heavy>() { return CRYPTONIGHT_HEAVY_MASK; } + +inline size_t cn_select_mask(xmrstak_algo algo) +{ + switch(algo) + { + case cryptonight: + return CRYPTONIGHT_MASK; + case cryptonight_lite: + return CRYPTONIGHT_LITE_MASK; + case cryptonight_monero: + return CRYPTONIGHT_MASK; + case cryptonight_heavy: + return CRYPTONIGHT_HEAVY_MASK; + default: + return 0; + } +} + +template<xmrstak_algo ALGO> +inline constexpr uint32_t cn_select_iter() { return 0; } + +template<> +inline constexpr uint32_t cn_select_iter<cryptonight>() { return CRYPTONIGHT_ITER; } + +template<> +inline constexpr uint32_t cn_select_iter<cryptonight_lite>() { return CRYPTONIGHT_LITE_ITER; } + +template<> +inline constexpr uint32_t cn_select_iter<cryptonight_monero>() { return CRYPTONIGHT_ITER; } -// define xmr settings -#define MONERO_MEMORY 2097152llu -#define MONERO_MASK 0x1FFFF0 -#define MONERO_ITER 0x80000 +template<> +inline constexpr uint32_t cn_select_iter<cryptonight_heavy>() { return CRYPTONIGHT_HEAVY_ITER; } +inline size_t cn_select_iter(xmrstak_algo algo) +{ + switch(algo) + { + case cryptonight: + return CRYPTONIGHT_ITER; + case cryptonight_lite: + return CRYPTONIGHT_LITE_ITER; + case cryptonight_monero: + return CRYPTONIGHT_ITER; + case cryptonight_heavy: + return CRYPTONIGHT_HEAVY_ITER; + default: + return 0; + } +} diff --git a/xmrstak/backend/miner_work.hpp b/xmrstak/backend/miner_work.hpp index 4bfe429..9e5a4e4 100644 --- a/xmrstak/backend/miner_work.hpp +++ b/xmrstak/backend/miner_work.hpp @@ -74,5 +74,11 @@ namespace xmrstak return *this; } + + uint8_t getVersion() const + { + return bWorkBlob[0]; + } + }; } // namepsace xmrstak diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp index 867a998..153e4e3 100644 --- a/xmrstak/backend/nvidia/minethd.cpp +++ b/xmrstak/backend/nvidia/minethd.cpp @@ -237,12 +237,20 @@ void minethd::work_main() uint64_t iCount = 0; cryptonight_ctx* cpu_ctx; cpu_ctx = cpu::minethd::minethd_alloc_ctx(); - cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, ::jconf::inst()->IsCurrencyMonero()); + auto miner_algo = ::jconf::inst()->GetMiningAlgo(); + cn_hash_fun hash_fun; + if(miner_algo == cryptonight_monero || miner_algo == cryptonight_heavy) + { + // start with cryptonight and switch later if fork version is reached + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight); + } + else + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); uint32_t iNonce; globalStates::inst().iConsumeCnt++; - bool mineMonero = strcmp_i(::jconf::inst()->GetCurrency(), "monero"); + uint8_t version = 0; while (bQuit == 0) { @@ -257,6 +265,16 @@ void minethd::work_main() std::this_thread::sleep_for(std::chrono::milliseconds(100)); consume_work(); + uint8_t new_version = oWork.getVersion(); + if(miner_algo == cryptonight_monero && version < 7 && new_version >= 7) + { + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_monero); + } + else if(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3) + { + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy); + } + version = new_version; continue; } @@ -281,11 +299,11 @@ void minethd::work_main() uint32_t foundNonce[10]; uint32_t foundCount; - cryptonight_extra_cpu_prepare(&ctx, iNonce); + cryptonight_extra_cpu_prepare(&ctx, iNonce, miner_algo, version); - cryptonight_core_cpu_hash(&ctx, mineMonero); + cryptonight_core_cpu_hash(&ctx, miner_algo, iNonce, version); - cryptonight_extra_cpu_final(&ctx, iNonce, oWork.iTarget, &foundCount, foundNonce); + cryptonight_extra_cpu_final(&ctx, iNonce, oWork.iTarget, &foundCount, foundNonce, miner_algo, version); for(size_t i = 0; i < foundCount; i++) { @@ -316,6 +334,16 @@ void minethd::work_main() } consume_work(); + uint8_t new_version = oWork.getVersion(); + if(miner_algo == cryptonight_monero && version < 7 && new_version >= 7) + { + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_monero); + } + else if(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3) + { + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy); + } + version = new_version; } } diff --git a/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp b/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp index afbdbaf..29a3523 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp +++ b/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp @@ -3,6 +3,9 @@ #include <stdint.h> #include <string> +#include "xmrstak/jconf.hpp" +#include "xmrstak/backend/cryptonight.hpp" + typedef struct { int device_id; const char *device_name; @@ -20,6 +23,7 @@ typedef struct { uint32_t *d_result_nonce; uint32_t *d_long_state; uint32_t *d_ctx_state; + uint32_t *d_ctx_state2; uint32_t *d_ctx_a; uint32_t *d_ctx_b; uint32_t *d_ctx_key1; @@ -41,8 +45,8 @@ int cuda_get_devicecount( int* deviceCount); int cuda_get_deviceinfo(nvid_ctx *ctx); int cryptonight_extra_cpu_init(nvid_ctx *ctx); void cryptonight_extra_cpu_set_data( nvid_ctx* ctx, const void *data, uint32_t len); -void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce); -void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce); +void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce, xmrstak_algo miner_algo, uint8_t version); +void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce,xmrstak_algo miner_algo, uint8_t version); } -void cryptonight_core_cpu_hash(nvid_ctx* ctx, bool mineMonero); +void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t startNonce, uint8_t version); diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu index cc97274..ede578f 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu @@ -6,6 +6,8 @@ #include <cuda.h> #include <cuda_runtime.h> +#include "xmrstak/jconf.hpp" + #ifdef _WIN32 #include <windows.h> extern "C" void compat_usleep(uint64_t waitTime) @@ -106,8 +108,18 @@ __device__ __forceinline__ void storeGlobal32( T* addr, T const & val ) #endif } -template<size_t ITERATIONS, uint32_t THREAD_SHIFT> -__global__ void cryptonight_core_gpu_phase1( int threads, int bfactor, int partidx, uint32_t * __restrict__ long_state, uint32_t * __restrict__ ctx_state, uint32_t * __restrict__ ctx_key1 ) +template< typename T > +__device__ __forceinline__ void storeGlobal64( T* addr, T const & val ) +{ +#if (__CUDA_ARCH__ < 700) + asm volatile( "st.global.cg.u64 [%0], %1;" : : "l"( addr ), "l"( val ) ); +#else + *addr = val; +#endif +} + +template<size_t ITERATIONS, uint32_t MEMORY> +__global__ void cryptonight_core_gpu_phase1( int threads, int bfactor, int partidx, uint32_t * __restrict__ long_state, uint32_t * __restrict__ ctx_state2, uint32_t * __restrict__ ctx_key1 ) { __shared__ uint32_t sharedMemory[1024]; @@ -117,7 +129,7 @@ __global__ void cryptonight_core_gpu_phase1( int threads, int bfactor, int parti const int thread = ( blockDim.x * blockIdx.x + threadIdx.x ) >> 3; const int sub = ( threadIdx.x & 7 ) << 2; - const int batchsize = ITERATIONS >> bfactor; + const int batchsize = MEMORY >> bfactor; const int start = partidx * batchsize; const int end = start + batchsize; @@ -131,18 +143,18 @@ __global__ void cryptonight_core_gpu_phase1( int threads, int bfactor, int parti if( partidx == 0 ) { // first round - MEMCPY8( text, ctx_state + thread * 50 + sub + 16, 2 ); + MEMCPY8( text, ctx_state2 + thread * 50 + sub + 16, 2 ); } else { // load previous text data - MEMCPY8( text, &long_state[( (uint64_t) thread << THREAD_SHIFT ) + sub + start - 32], 2 ); + MEMCPY8( text, &long_state[( (uint64_t) thread * MEMORY ) + sub + start - 32], 2 ); } __syncthreads( ); for ( int i = start; i < end; i += 32 ) { cn_aes_pseudo_round_mut( sharedMemory, text, key ); - MEMCPY8(&long_state[((uint64_t) thread << THREAD_SHIFT) + (sub + i)], text, 2); + MEMCPY8(&long_state[((uint64_t) thread * MEMORY) + (sub + i)], text, 2); } } @@ -157,33 +169,37 @@ __forceinline__ __device__ void unusedVar( const T& ) * - this method can be used with all compute architectures * - for <sm_30 shared memory is needed * + * group_n - must be a power of 2! + * * @param ptr pointer to shared memory, size must be `threadIdx.x * sizeof(uint32_t)` * value can be NULL for compute architecture >=sm_30 - * @param sub thread number within the group, range [0;4) + * @param sub thread number within the group, range [0:group_n] * @param value value to share with other threads within the group - * @param src thread number within the group from where the data is read, range [0;4) + * @param src thread number within the group from where the data is read, range [0:group_n] */ +template<size_t group_n> __forceinline__ __device__ uint32_t shuffle(volatile uint32_t* ptr,const uint32_t sub,const int val,const uint32_t src) { #if( __CUDA_ARCH__ < 300 ) ptr[sub] = val; - return ptr[src&3]; + return ptr[src & (group_n-1)]; #else unusedVar( ptr ); unusedVar( sub ); # if(__CUDACC_VER_MAJOR__ >= 9) - return __shfl_sync(0xFFFFFFFF, val, src, 4 ); + return __shfl_sync(0xFFFFFFFF, val, src, group_n ); # else - return __shfl( val, src, 4 ); + return __shfl( val, src, group_n ); # endif #endif } -template<size_t ITERATIONS, uint32_t THREAD_SHIFT, uint32_t MASK> +template<size_t ITERATIONS, uint32_t MEMORY, uint32_t MASK, xmrstak_algo ALGO> #ifdef XMR_STAK_THREADS __launch_bounds__( XMR_STAK_THREADS * 4 ) #endif -__global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int partidx, uint32_t * d_long_state, uint32_t * d_ctx_a, uint32_t * d_ctx_b ) +__global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int partidx, uint32_t * d_long_state, uint32_t * d_ctx_a, uint32_t * d_ctx_b, uint32_t * d_ctx_state, + uint32_t startNonce, uint32_t * __restrict__ d_input ) { __shared__ uint32_t sharedMemory[1024]; @@ -192,6 +208,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti __syncthreads( ); const int thread = ( blockDim.x * blockIdx.x + threadIdx.x ) >> 2; + const uint32_t nonce = startNonce + thread; const int sub = threadIdx.x & 3; const int sub2 = sub & 2; @@ -205,30 +222,48 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti return; int i, k; - uint32_t j; + uint32_t j; const int batchsize = (ITERATIONS * 2) >> ( 2 + bfactor ); const int start = partidx * batchsize; const int end = start + batchsize; - uint32_t * long_state = &d_long_state[(IndexType) thread << THREAD_SHIFT]; - uint32_t * ctx_a = d_ctx_a + thread * 4; - uint32_t * ctx_b = d_ctx_b + thread * 4; - uint32_t a, d[2]; + uint32_t * long_state = &d_long_state[(IndexType) thread * MEMORY]; + uint32_t a, d[2], idx0; uint32_t t1[2], t2[2], res; - a = ctx_a[sub]; - d[1] = ctx_b[sub]; + uint32_t tweak1_2[2]; + if (ALGO == cryptonight_monero) + { + uint32_t * state = d_ctx_state + thread * 50; + tweak1_2[0] = (d_input[8] >> 24) | (d_input[9] << 8); + tweak1_2[0] ^= state[48]; + tweak1_2[1] = nonce; + tweak1_2[1] ^= state[49]; + } + + a = (d_ctx_a + thread * 4)[sub]; + idx0 = shuffle<4>(sPtr,sub, a, 0); + if(ALGO == cryptonight_heavy) + { + if(partidx != 0) + { + // state is stored after all ctx_b states + idx0 = *(d_ctx_b + threads * 4 + thread); + } + } + d[1] = (d_ctx_b + thread * 4)[sub]; + #pragma unroll 2 for ( i = start; i < end; ++i ) { #pragma unroll 2 for ( int x = 0; x < 2; ++x ) { - j = ( ( shuffle(sPtr,sub, a, 0) & MASK ) >> 2 ) + sub; + j = ( ( idx0 & MASK ) >> 2 ) + sub; const uint32_t x_0 = loadGlobal32<uint32_t>( long_state + j ); - const uint32_t x_1 = shuffle(sPtr,sub, x_0, sub + 1); - const uint32_t x_2 = shuffle(sPtr,sub, x_0, sub + 2); - const uint32_t x_3 = shuffle(sPtr,sub, x_0, sub + 3); + const uint32_t x_1 = shuffle<4>(sPtr,sub, x_0, sub + 1); + const uint32_t x_2 = shuffle<4>(sPtr,sub, x_0, sub + 2); + const uint32_t x_3 = shuffle<4>(sPtr,sub, x_0, sub + 3); d[x] = a ^ t_fn0( x_0 & 0xff ) ^ t_fn1( (x_1 >> 8) & 0xff ) ^ @@ -237,41 +272,74 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti //XOR_BLOCKS_DST(c, b, &long_state[j]); - t1[0] = shuffle(sPtr,sub, d[x], 0); - //long_state[j] = d[0] ^ d[1]; - storeGlobal32( long_state + j, d[0] ^ d[1] ); - + t1[0] = shuffle<4>(sPtr,sub, d[x], 0); + + const uint32_t z = d[0] ^ d[1]; + if(ALGO == cryptonight_monero) + { + const uint32_t table = 0x75310U; + const uint32_t index = ((z >> 26) & 12) | ((z >> 23) & 2); + const uint32_t fork_7 = z ^ ((table >> index) & 0x30U) << 24; + storeGlobal32( long_state + j, sub == 2 ? fork_7 : z ); + } + else + storeGlobal32( long_state + j, z ); + //MUL_SUM_XOR_DST(c, a, &long_state[((uint32_t *)c)[0] & MASK]); j = ( ( *t1 & MASK ) >> 2 ) + sub; uint32_t yy[2]; *( (uint64_t*) yy ) = loadGlobal64<uint64_t>( ( (uint64_t *) long_state )+( j >> 1 ) ); uint32_t zz[2]; - zz[0] = shuffle(sPtr,sub, yy[0], 0); - zz[1] = shuffle(sPtr,sub, yy[1], 0); + zz[0] = shuffle<4>(sPtr,sub, yy[0], 0); + zz[1] = shuffle<4>(sPtr,sub, yy[1], 0); - t1[1] = shuffle(sPtr,sub, d[x], 1); + t1[1] = shuffle<4>(sPtr,sub, d[x], 1); #pragma unroll for ( k = 0; k < 2; k++ ) - t2[k] = shuffle(sPtr,sub, a, k + sub2); + t2[k] = shuffle<4>(sPtr,sub, a, k + sub2); *( (uint64_t *) t2 ) += sub2 ? ( *( (uint64_t *) t1 ) * *( (uint64_t*) zz ) ) : __umul64hi( *( (uint64_t *) t1 ), *( (uint64_t*) zz ) ); res = *( (uint64_t *) t2 ) >> ( sub & 1 ? 32 : 0 ); - storeGlobal32( long_state + j, res ); + + if(ALGO == cryptonight_monero) + { + const uint32_t tweaked_res = tweak1_2[sub & 1] ^ res; + const uint32_t long_state_update = sub2 ? tweaked_res : res; + storeGlobal32( long_state + j, long_state_update ); + } + else + storeGlobal32( long_state + j, res ); + a = ( sub & 1 ? yy[1] : yy[0] ) ^ res; + idx0 = shuffle<4>(sPtr,sub, a, 0); + if(ALGO == cryptonight_heavy) + { + int64_t n = loadGlobal64<uint64_t>( ( (uint64_t *) long_state ) + (( idx0 & MASK ) >> 3)); + int32_t d = loadGlobal32<uint32_t>( (uint32_t*)(( (uint64_t *) long_state ) + (( idx0 & MASK) >> 3) + 1u )); + int64_t q = n / (d | 0x5); + + if(sub&1) + storeGlobal64<uint64_t>( ( (uint64_t *) long_state ) + (( idx0 & MASK ) >> 3), n ^ q ); + + idx0 = d ^ q; + } } } if ( bfactor > 0 ) { - ctx_a[sub] = a; - ctx_b[sub] = d[1]; + (d_ctx_a + thread * 4)[sub] = a; + (d_ctx_b + thread * 4)[sub] = d[1]; + if(ALGO == cryptonight_heavy) + if(sub&1) + *(d_ctx_b + threads * 4 + thread) = idx0; } } -template<size_t ITERATIONS, uint32_t THREAD_SHIFT> +template<size_t ITERATIONS, uint32_t MEMORY, xmrstak_algo ALGO> __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int partidx, const uint32_t * __restrict__ long_state, uint32_t * __restrict__ d_ctx_state, uint32_t * __restrict__ d_ctx_key2 ) { __shared__ uint32_t sharedMemory[1024]; @@ -280,9 +348,10 @@ __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int parti __syncthreads( ); int thread = ( blockDim.x * blockIdx.x + threadIdx.x ) >> 3; - int sub = ( threadIdx.x & 7 ) << 2; + int subv = ( threadIdx.x & 7 ); + int sub = subv << 2; - const int batchsize = ITERATIONS >> bfactor; + const int batchsize = MEMORY >> bfactor; const int start = partidx * batchsize; const int end = start + batchsize; @@ -294,20 +363,53 @@ __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int parti MEMCPY8( text, d_ctx_state + thread * 50 + sub + 16, 2 ); __syncthreads( ); + +#if( __CUDA_ARCH__ < 300 ) + extern __shared__ uint32_t shuffleMem[]; + volatile uint32_t* sPtr = (volatile uint32_t*)(shuffleMem + (threadIdx.x& 0xFFFFFFFC)); +#else + volatile uint32_t* sPtr = NULL; +#endif + for ( int i = start; i < end; i += 32 ) { #pragma unroll for ( int j = 0; j < 4; ++j ) - text[j] ^= long_state[((IndexType) thread << THREAD_SHIFT) + (sub + i + j)]; + text[j] ^= long_state[((IndexType) thread * MEMORY) + ( sub + i + j)]; cn_aes_pseudo_round_mut( sharedMemory, text, key ); + + if(ALGO == cryptonight_heavy) + { +#pragma unroll + for ( int j = 0; j < 4; ++j ) + text[j] ^= shuffle<8>(sPtr, subv, text[j], (subv+1)&7); + } + } + + if(ALGO == cryptonight_heavy) + { + __syncthreads( ); + + for ( int i = start; i < end; i += 32 ) + { +#pragma unroll + for ( int j = 0; j < 4; ++j ) + text[j] ^= long_state[((IndexType) thread * MEMORY) + ( sub + i + j)]; + + cn_aes_pseudo_round_mut( sharedMemory, text, key ); + +#pragma unroll + for ( int j = 0; j < 4; ++j ) + text[j] ^= shuffle<8>(sPtr, subv, text[j], (subv+1)&7); + } } MEMCPY8( d_ctx_state + thread * 50 + sub + 16, text, 2 ); } -template<size_t ITERATIONS, uint32_t MASK, uint32_t THREAD_SHIFT> -void cryptonight_core_gpu_hash(nvid_ctx* ctx) +template<size_t ITERATIONS, uint32_t MASK, uint32_t MEMORY, xmrstak_algo ALGO> +void cryptonight_core_gpu_hash(nvid_ctx* ctx, uint32_t nonce) { dim3 grid( ctx->device_blocks ); dim3 block( ctx->device_threads ); @@ -329,9 +431,11 @@ void cryptonight_core_gpu_hash(nvid_ctx* ctx) for ( int i = 0; i < partcountOneThree; i++ ) { - CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase1<ITERATIONS,THREAD_SHIFT><<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads, + CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase1<ITERATIONS,MEMORY><<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads, bfactorOneThree, i, - ctx->d_long_state, ctx->d_ctx_state, ctx->d_ctx_key1 )); + ctx->d_long_state, + (ALGO == cryptonight_heavy ? ctx->d_ctx_state2 : ctx->d_ctx_state), + ctx->d_ctx_key1 )); if ( partcount > 1 && ctx->device_bsleep > 0) compat_usleep( ctx->device_bsleep ); } @@ -342,7 +446,7 @@ void cryptonight_core_gpu_hash(nvid_ctx* ctx) CUDA_CHECK_MSG_KERNEL( ctx->device_id, "\n**suggestion: Try to increase the value of the attribute 'bfactor' or \nreduce 'threads' in the NVIDIA config file.**", - cryptonight_core_gpu_phase2<ITERATIONS,THREAD_SHIFT,MASK><<< + cryptonight_core_gpu_phase2<ITERATIONS,MEMORY,MASK,ALGO><<< grid, block4, block4.x * sizeof(uint32_t) * static_cast< int >( ctx->device_arch[0] < 3 ) @@ -352,7 +456,10 @@ void cryptonight_core_gpu_hash(nvid_ctx* ctx) i, ctx->d_long_state, ctx->d_ctx_a, - ctx->d_ctx_b + ctx->d_ctx_b, + ctx->d_ctx_state, + nonce, + ctx->d_input ) ); @@ -361,25 +468,39 @@ void cryptonight_core_gpu_hash(nvid_ctx* ctx) for ( int i = 0; i < partcountOneThree; i++ ) { - CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase3<ITERATIONS,THREAD_SHIFT><<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads, + CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase3<ITERATIONS,MEMORY, ALGO><<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads, bfactorOneThree, i, ctx->d_long_state, ctx->d_ctx_state, ctx->d_ctx_key2 )); } } -void cryptonight_core_cpu_hash(nvid_ctx* ctx, bool mineMonero) +void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t startNonce, uint8_t version) { -#ifndef CONF_NO_MONERO - if(mineMonero) + + if(miner_algo == cryptonight_monero) { - cryptonight_core_gpu_hash<MONERO_ITER, MONERO_MASK, 19u>(ctx); + if(version >= 7) + cryptonight_core_gpu_hash<CRYPTONIGHT_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight_monero>(ctx, startNonce); + else + cryptonight_core_gpu_hash<CRYPTONIGHT_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight>(ctx, startNonce); } -#endif -#ifndef CONF_NO_AEON - if(!mineMonero) + else if(miner_algo == cryptonight_heavy) { - cryptonight_core_gpu_hash<AEON_ITER, AEON_MASK, 18u>(ctx); + if(version >= 3) + cryptonight_core_gpu_hash<CRYPTONIGHT_HEAVY_ITER, CRYPTONIGHT_HEAVY_MASK, CRYPTONIGHT_HEAVY_MEMORY/4, cryptonight_heavy>(ctx, startNonce); + else + { + cryptonight_core_gpu_hash<CRYPTONIGHT_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight>(ctx, startNonce); + } } -#endif + else if(miner_algo == cryptonight) + { + cryptonight_core_gpu_hash<CRYPTONIGHT_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight>(ctx, startNonce); + } + else if(miner_algo == cryptonight_lite) + { + cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_lite>(ctx, startNonce); + } + } diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu index 92259db..2f08a1a 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu @@ -28,6 +28,7 @@ uint64_t keccakf_rndc[24] ={ typedef unsigned char BitSequence; typedef unsigned long long DataLength; +#include "xmrstak/backend/cryptonight.hpp" #include "cryptonight.hpp" #include "cuda_extra.hpp" #include "cuda_keccak.hpp" @@ -36,6 +37,7 @@ typedef unsigned long long DataLength; #include "cuda_jh.hpp" #include "cuda_skein.hpp" #include "cuda_device.hpp" +#include "cuda_aes.hpp" __constant__ uint8_t d_sub_byte[16][16] ={ {0x63, 0x7c, 0x77, 0x7b, 0xf2, 0x6b, 0x6f, 0xc5, 0x30, 0x01, 0x67, 0x2b, 0xfe, 0xd7, 0xab, 0x76 }, @@ -90,10 +92,33 @@ __device__ __forceinline__ void cryptonight_aes_set_key( uint32_t * __restrict__ } } -__global__ void cryptonight_extra_gpu_prepare( int threads, uint32_t * __restrict__ d_input, uint32_t len, uint32_t startNonce, uint32_t * __restrict__ d_ctx_state, uint32_t * __restrict__ d_ctx_a, uint32_t * __restrict__ d_ctx_b, uint32_t * __restrict__ d_ctx_key1, uint32_t * __restrict__ d_ctx_key2 ) +__device__ __forceinline__ void mix_and_propagate( uint32_t* state ) +{ + uint32_t tmp0[4]; + for(size_t x = 0; x < 4; ++x) + tmp0[x] = (state)[x]; + + // set destination [0,6] + for(size_t t = 0; t < 7; ++t) + for(size_t x = 0; x < 4; ++x) + (state + 4 * t)[x] = (state + 4 * t)[x] ^ (state + 4 * (t + 1))[x]; + + // set destination 7 + for(size_t x = 0; x < 4; ++x) + (state + 4 * 7)[x] = (state + 4 * 7)[x] ^ tmp0[x]; +} + +template<xmrstak_algo ALGO> +__global__ void cryptonight_extra_gpu_prepare( int threads, uint32_t * __restrict__ d_input, uint32_t len, uint32_t startNonce, uint32_t * __restrict__ d_ctx_state, uint32_t * __restrict__ d_ctx_state2, uint32_t * __restrict__ d_ctx_a, uint32_t * __restrict__ d_ctx_b, uint32_t * __restrict__ d_ctx_key1, uint32_t * __restrict__ d_ctx_key2 ) { int thread = ( blockDim.x * blockIdx.x + threadIdx.x ); + __shared__ uint32_t sharedMemory[1024]; + if(ALGO == cryptonight_heavy) + { + cn_aes_gpu_init( sharedMemory ); + __syncthreads( ); + } if ( thread >= threads ) return; @@ -113,20 +138,45 @@ __global__ void cryptonight_extra_gpu_prepare( int threads, uint32_t * __restric cn_keccak( (uint8_t *) input, len, (uint8_t *) ctx_state ); cryptonight_aes_set_key( ctx_key1, ctx_state ); cryptonight_aes_set_key( ctx_key2, ctx_state + 8 ); + XOR_BLOCKS_DST( ctx_state, ctx_state + 8, ctx_a ); XOR_BLOCKS_DST( ctx_state + 4, ctx_state + 12, ctx_b ); - - memcpy( d_ctx_state + thread * 50, ctx_state, 50 * 4 ); memcpy( d_ctx_a + thread * 4, ctx_a, 4 * 4 ); memcpy( d_ctx_b + thread * 4, ctx_b, 4 * 4 ); + memcpy( d_ctx_key1 + thread * 40, ctx_key1, 40 * 4 ); memcpy( d_ctx_key2 + thread * 40, ctx_key2, 40 * 4 ); + memcpy( d_ctx_state + thread * 50, ctx_state, 50 * 4 ); + + if(ALGO == cryptonight_heavy) + { + + for(int i=0; i < 16; i++) + { + for(size_t t = 4; t < 12; ++t) + { + cn_aes_pseudo_round_mut( sharedMemory, ctx_state + 4u * t, ctx_key1 ); + } + // scipt first 4 * 128bit blocks = 4 * 4 uint32_t values + mix_and_propagate(ctx_state + 4 * 4); + } + // double buffer to move manipulated state into phase1 + memcpy( d_ctx_state2 + thread * 50, ctx_state, 50 * 4 ); + } } -__global__ void cryptonight_extra_gpu_final( int threads, uint64_t target, uint32_t* __restrict__ d_res_count, uint32_t * __restrict__ d_res_nonce, uint32_t * __restrict__ d_ctx_state ) +template<xmrstak_algo ALGO> +__global__ void cryptonight_extra_gpu_final( int threads, uint64_t target, uint32_t* __restrict__ d_res_count, uint32_t * __restrict__ d_res_nonce, uint32_t * __restrict__ d_ctx_state,uint32_t * __restrict__ d_ctx_key2 ) { const int thread = blockDim.x * blockIdx.x + threadIdx.x; + __shared__ uint32_t sharedMemory[1024]; + + if(ALGO == cryptonight_heavy) + { + cn_aes_gpu_init( sharedMemory ); + __syncthreads( ); + } if ( thread >= threads ) return; @@ -134,11 +184,28 @@ __global__ void cryptonight_extra_gpu_final( int threads, uint64_t target, uint3 uint32_t * __restrict__ ctx_state = d_ctx_state + thread * 50; uint64_t hash[4]; uint32_t state[50]; - -#pragma unroll + + #pragma unroll for ( i = 0; i < 50; i++ ) state[i] = ctx_state[i]; + if(ALGO == cryptonight_heavy) + { + uint32_t key[40]; + + // load keys + MEMCPY8( key, d_ctx_key2 + thread * 40, 20 ); + + for(int i=0; i < 16; i++) + { + for(size_t t = 4; t < 12; ++t) + { + cn_aes_pseudo_round_mut( sharedMemory, state + 4u * t, key ); + } + // scipt first 4 * 128bit blocks = 4 * 4 uint32_t values + mix_and_propagate(state + 4 * 4); + } + } cn_keccakf2( (uint64_t *) state ); switch ( ( (uint8_t *) state )[0] & 0x03 ) @@ -212,23 +279,26 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx) if(gpuArch < 70) CUDA_CHECK(ctx->device_id, cudaDeviceSetCacheConfig(cudaFuncCachePreferL1)); - size_t hashMemSize; - if(::jconf::inst()->IsCurrencyMonero()) - { - hashMemSize = MONERO_MEMORY; - } - else - { - hashMemSize = AEON_MEMORY; - } + size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo()); size_t wsize = ctx->device_blocks * ctx->device_threads; CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_state, 50 * sizeof(uint32_t) * wsize)); + size_t ctx_b_size = 4 * sizeof(uint32_t) * wsize; + if(cryptonight_heavy == ::jconf::inst()->GetMiningAlgo()) + { + // extent ctx_b to hold the state of idx0 + ctx_b_size += sizeof(uint32_t) * wsize; + // create a double buffer for the state to exchange the mixed state to phase1 + CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_state2, 50 * sizeof(uint32_t) * wsize)); + } + else + ctx->d_ctx_state2 = ctx->d_ctx_state; + CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_key1, 40 * sizeof(uint32_t) * wsize)); CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_key2, 40 * sizeof(uint32_t) * wsize)); CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_text, 32 * sizeof(uint32_t) * wsize)); CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_a, 4 * sizeof(uint32_t) * wsize)); - CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_b, 4 * sizeof(uint32_t) * wsize)); + CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_b, ctx_b_size)); CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_input, 21 * sizeof (uint32_t ) )); CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_result_count, sizeof (uint32_t ) )); CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_result_nonce, 10 * sizeof (uint32_t ) )); @@ -239,7 +309,7 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx) return 1; } -extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce) +extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce, xmrstak_algo miner_algo, uint8_t version) { int threadsperblock = 128; uint32_t wsize = ctx->device_blocks * ctx->device_threads; @@ -247,11 +317,22 @@ extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce dim3 grid( ( wsize + threadsperblock - 1 ) / threadsperblock ); dim3 block( threadsperblock ); - CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_extra_gpu_prepare<<<grid, block >>>( wsize, ctx->d_input, ctx->inputlen, startNonce, - ctx->d_ctx_state, ctx->d_ctx_a, ctx->d_ctx_b, ctx->d_ctx_key1, ctx->d_ctx_key2 )); + if(miner_algo == cryptonight_heavy && version >= 3) + { + CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_extra_gpu_prepare<cryptonight_heavy><<<grid, block >>>( wsize, ctx->d_input, ctx->inputlen, startNonce, + ctx->d_ctx_state,ctx->d_ctx_state2, ctx->d_ctx_a, ctx->d_ctx_b, ctx->d_ctx_key1, ctx->d_ctx_key2 )); + } + else + { + /* pass two times d_ctx_state because the second state is used later in phase1, + * the first is used than in phase3 + */ + CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_extra_gpu_prepare<invalid_algo><<<grid, block >>>( wsize, ctx->d_input, ctx->inputlen, startNonce, + ctx->d_ctx_state, ctx->d_ctx_state, ctx->d_ctx_a, ctx->d_ctx_b, ctx->d_ctx_key1, ctx->d_ctx_key2 )); + } } -extern "C" void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce) +extern "C" void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce,xmrstak_algo miner_algo, uint8_t version) { int threadsperblock = 128; uint32_t wsize = ctx->device_blocks * ctx->device_threads; @@ -262,11 +343,23 @@ extern "C" void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, CUDA_CHECK(ctx->device_id, cudaMemset( ctx->d_result_nonce, 0xFF, 10 * sizeof (uint32_t ) )); CUDA_CHECK(ctx->device_id, cudaMemset( ctx->d_result_count, 0, sizeof (uint32_t ) )); - CUDA_CHECK_MSG_KERNEL( - ctx->device_id, - "\n**suggestion: Try to increase the value of the attribute 'bfactor' in the NVIDIA config file.**", - cryptonight_extra_gpu_final<<<grid, block >>>( wsize, target, ctx->d_result_count, ctx->d_result_nonce, ctx->d_ctx_state ) - ); + if(miner_algo == cryptonight_heavy && version >= 3) + { + CUDA_CHECK_MSG_KERNEL( + ctx->device_id, + "\n**suggestion: Try to increase the value of the attribute 'bfactor' in the NVIDIA config file.**", + cryptonight_extra_gpu_final<cryptonight_heavy><<<grid, block >>>( wsize, target, ctx->d_result_count, ctx->d_result_nonce, ctx->d_ctx_state,ctx->d_ctx_key2 ) + ); + } + else + { + // fallback for all other algorithms + CUDA_CHECK_MSG_KERNEL( + ctx->device_id, + "\n**suggestion: Try to increase the value of the attribute 'bfactor' in the NVIDIA config file.**", + cryptonight_extra_gpu_final<invalid_algo><<<grid, block >>>( wsize, target, ctx->d_result_count, ctx->d_result_nonce, ctx->d_ctx_state,ctx->d_ctx_key2 ) + ); + } CUDA_CHECK(ctx->device_id, cudaMemcpy( rescount, ctx->d_result_count, sizeof (uint32_t ), cudaMemcpyDeviceToHost )); CUDA_CHECK(ctx->device_id, cudaMemcpy( resnonce, ctx->d_result_nonce, 10 * sizeof (uint32_t ), cudaMemcpyDeviceToHost )); @@ -482,15 +575,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) ctx->total_device_memory = totalMemory; ctx->free_device_memory = freeMemory; - size_t hashMemSize; - if(::jconf::inst()->IsCurrencyMonero()) - { - hashMemSize = MONERO_MEMORY; - } - else - { - hashMemSize = AEON_MEMORY; - } + size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo()); #ifdef WIN32 /* We use in windows bfactor (split slow kernel into smaller parts) to avoid @@ -520,6 +605,9 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) // up to 16kibyte extra memory is used per thread for some kernel (lmem/local memory) // 680bytes are extra meta data memory per hash size_t perThread = hashMemSize + 16192u + 680u; + if(cryptonight_heavy == ::jconf::inst()->GetMiningAlgo()) + perThread += 50 * 4; // state double buffer + size_t max_intensity = limitedMemory / perThread; ctx->device_threads = max_intensity / ctx->device_blocks; // use only odd number of threads diff --git a/xmrstak/cli/cli-miner.cpp b/xmrstak/cli/cli-miner.cpp index 9053844..e2d50f2 100644 --- a/xmrstak/cli/cli-miner.cpp +++ b/xmrstak/cli/cli-miner.cpp @@ -55,7 +55,7 @@ # include "xmrstak/misc/uac.hpp" #endif // _WIN32 -void do_benchmark(); +int do_benchmark(int block_version); void help() { @@ -64,47 +64,50 @@ void help() cout<<"Usage: "<<params::inst().binaryName<<" [OPTION]..."<<endl; cout<<" "<<endl; - cout<<" -h, --help show this help"<<endl; - cout<<" -v, --version show version number"<<endl; - cout<<" -V, --version-long show long version number"<<endl; - cout<<" -c, --config FILE common miner configuration file"<<endl; + cout<<" -h, --help show this help"<<endl; + cout<<" -v, --version show version number"<<endl; + cout<<" -V, --version-long show long version number"<<endl; + cout<<" -c, --config FILE common miner configuration file"<<endl; + cout<<" -C, --poolconf FILE pool configuration file"<<endl; #ifdef _WIN32 - cout<<" --noUAC disable the UAC dialog"<<endl; -#endif -#if (!defined(CONF_NO_AEON)) && (!defined(CONF_NO_MONERO)) - cout<<" --currency NAME currency to mine: monero or aeon"<<endl; + cout<<" --noUAC disable the UAC dialog"<<endl; #endif + cout<<" --currency NAME currency to mine:"<<endl; + cout<<" --benchmark BLOCKVERSION ONLY do a 60-second benchmark and exit"<<endl; #ifndef CONF_NO_CPU - cout<<" --noCPU disable the CPU miner backend"<<endl; - cout<<" --cpu FILE CPU backend miner config file"<<endl; + cout<<" --noCPU disable the CPU miner backend"<<endl; + cout<<" --cpu FILE CPU backend miner config file"<<endl; #endif #ifndef CONF_NO_OPENCL - cout<<" --noAMD disable the AMD miner backend"<<endl; - cout<<" --amd FILE AMD backend miner config file"<<endl; + cout<<" --noAMD disable the AMD miner backend"<<endl; + cout<<" --amd FILE AMD backend miner config file"<<endl; #endif #ifndef CONF_NO_CUDA - cout<<" --noNVIDIA disable the NVIDIA miner backend"<<endl; - cout<<" --nvidia FILE NVIDIA backend miner config file"<<endl; + 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; + 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; + 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<< endl; #ifdef _WIN32 cout<<"Environment variables:\n"<<endl; - cout<<" XMRSTAK_NOWAIT disable the dialog `Press any key to exit."<<std::endl; - cout<<" for non UAC execution"<<endl; - cout<<" \n"<<endl; + cout<<" XMRSTAK_NOWAIT disable the dialog `Press any key to exit."<<std::endl; + cout<<" for non UAC execution"<<endl; + cout<< endl; #endif + std::string algos; + jconf::GetAlgoList(algos); + cout<< "Supported coin opitons: " << endl << algos << endl; cout<< "Version: " << get_version_str_short() << endl; cout<<"Brought to by fireice_uk and psychocrypt under GPLv3."<<endl; } @@ -133,10 +136,7 @@ std::string get_multipool_entry(bool& final) std::cout<<std::endl<<"- Next Pool:"<<std::endl<<std::endl; std::string pool; - if(xmrstak::params::inst().currency == "monero") - std::cout<<"- Pool address: e.g. pool.usxmrpool.com:3333"<<std::endl; - else - std::cout<<"- Pool address: e.g. mine.aeon-pool.com:5555"<<std::endl; + std::cout<<"- Pool address: e.g. " << jconf::GetDefaultPool(xmrstak::params::inst().currency.c_str()) << std::endl; std::cin >> pool; std::string userName; @@ -149,7 +149,6 @@ std::string get_multipool_entry(bool& final) 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); @@ -185,61 +184,36 @@ inline void prompt_once(bool& prompted) } } -void do_guided_config() +void do_guided_pool_config() { using namespace xmrstak; // load the template of the backend config into a char variable const char *tpl = - #include "../config.tpl" + #include "../pools.tpl" ; configEditor configTpl{}; configTpl.set(std::string(tpl)); bool prompted = false; - + auto& currency = params::inst().currency; - if(currency.empty()) + if(currency.empty() || !jconf::IsOnAlgoList(currency)) { prompt_once(prompted); std::string tmp; -#if defined(CONF_NO_AEON) - tmp = "monero"; -#elif defined(CONF_NO_MONERO) - tmp = "aeon"; -#endif - while(tmp != "monero" && tmp != "aeon") + while(tmp.empty() || !jconf::IsOnAlgoList(tmp)) { - std::cout<<"- Currency: 'monero' or 'aeon'"<<std::endl; + std::string list; + jconf::GetAlgoList(list); + std::cout << "- Please enter the currency that you want to mine: "<<std::endl; + std::cout << list << std::endl; std::cin >> tmp; - std::transform(tmp.begin(), tmp.end(), tmp.begin(), ::tolower); } 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()) @@ -247,10 +221,7 @@ void do_guided_config() prompt_once(prompted); userSetPool = false; - if(currency == "monero") - std::cout<<"- Pool address: e.g. pool.usxmrpool.com:3333"<<std::endl; - else - std::cout<<"- Pool address: e.g. mine.aeon-pool.com:5555"<<std::endl; + std::cout<<"- Pool address: e.g. " << jconf::GetDefaultPool(xmrstak::params::inst().currency.c_str()) << std::endl; std::cin >> pool; } @@ -263,6 +234,7 @@ void do_guided_config() std::cin >> userName; } + bool stdin_flushed = false; auto& passwd = params::inst().poolPasswd; if(passwd.empty() && !params::inst().userSetPwd) { @@ -270,6 +242,8 @@ void do_guided_config() // clear everything from stdin to allow an empty password std::cin.clear(); std::cin.ignore(INT_MAX,'\n'); + stdin_flushed = true; + std::cout<<"- Password (mostly empty or x):"<<std::endl; getline(std::cin, passwd); } @@ -279,8 +253,12 @@ void do_guided_config() { prompt_once(prompted); - // clear everything from stdin to allow an empty rigid - std::cin.clear(); std::cin.ignore(INT_MAX,'\n'); + if(!stdin_flushed) + { + // 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); } @@ -346,8 +324,49 @@ void do_guided_config() while(!final); } - configTpl.replace("POOLCONF", pool_table); configTpl.replace("CURRENCY", currency); + configTpl.replace("POOLCONF", pool_table); + configTpl.write(params::inst().configFilePools); + std::cout<<"Pool configuration stored in file '"<<params::inst().configFilePools<<"'"<<std::endl; +} + +void do_guided_config() +{ + using namespace xmrstak; + + // load the template of the backend config into a char variable + const char *tpl = + #include "../config.tpl" + ; + + configEditor configTpl{}; + configTpl.set(std::string(tpl)); + bool prompted = false; + + 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 + prompt_once(prompted); + + 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 + } + 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; @@ -574,6 +593,17 @@ int main(int argc, char *argv[]) } params::inst().configFile = argv[i]; } + else if(opName.compare("-C") == 0 || opName.compare("--poolconf") == 0) + { + ++i; + if( i >=argc ) + { + printer::inst()->print_msg(L0, "No argument for parameter '-C/--poolconf' given"); + win_exit(); + return 1; + } + params::inst().configFilePools = argv[i]; + } else if(opName.compare("-i") == 0 || opName.compare("--httpd") == 0) { ++i; @@ -600,6 +630,25 @@ int main(int argc, char *argv[]) { params::inst().allowUAC = false; } + else if(opName.compare("--benchmark") == 0) + { + ++i; + if( i >= argc ) + { + printer::inst()->print_msg(L0, "No argument for parameter '--benchmark' given"); + win_exit(); + return 1; + } + char* block_version = nullptr; + long int bversion = strtol(argv[i], &block_version, 10); + + if(bversion < 0 || bversion >= 256) + { + printer::inst()->print_msg(L0, "Benchmark block version must be in the range [0,255]"); + return 1; + } + params::inst().benchmark_block_version = bversion; + } else { printer::inst()->print_msg(L0, "Parameter unknown '%s'",argv[i]); @@ -612,7 +661,10 @@ int main(int argc, char *argv[]) if(!configEditor::file_exist(params::inst().configFile)) do_guided_config(); - if(!jconf::inst()->parse_config(params::inst().configFile.c_str())) + if(!configEditor::file_exist(params::inst().configFilePools)) + do_guided_pool_config(); + + if(!jconf::inst()->parse_config(params::inst().configFile.c_str(), params::inst().configFilePools.c_str())) { win_exit(); return 1; @@ -670,11 +722,14 @@ int main(int argc, char *argv[]) printer::inst()->print_str("'r' - results\n"); printer::inst()->print_str("'c' - connection\n"); printer::inst()->print_str("-------------------------------------------------------------------\n"); - if(::jconf::inst()->IsCurrencyMonero()) - printer::inst()->print_msg(L0,"Start mining: MONERO"); - else - printer::inst()->print_msg(L0,"Start mining: AEON"); + printer::inst()->print_msg(L0, "Mining coin: %s", jconf::inst()->GetMiningCoin().c_str()); + if(params::inst().benchmark_block_version >= 0) + { + printer::inst()->print_str("!!!! Doing only a benchmark and exiting. To mine, remove the '--benchmark' option. !!!!\n"); + return do_benchmark(params::inst().benchmark_block_version); + } + executor::inst()->ex_start(jconf::inst()->DaemonMode()); uint64_t lastTime = get_timestamp_ms(); @@ -709,23 +764,31 @@ int main(int argc, char *argv[]) return 0; } -void do_benchmark() +int do_benchmark(int block_version) { using namespace std::chrono; std::vector<xmrstak::iBackend*>* pvThreads; - printer::inst()->print_msg(L0, "Running a 60 second benchmark..."); + printer::inst()->print_msg(L0, "Prepare benchmark for block version %d", block_version); + + uint8_t work[112]; + memset(work,0,112); + work[0] = static_cast<uint8_t>(block_version); - uint8_t work[76] = {0}; - xmrstak::miner_work oWork = xmrstak::miner_work("", work, sizeof(work), 0, false, 0); + xmrstak::pool_data dat; + + xmrstak::miner_work oWork = xmrstak::miner_work(); pvThreads = xmrstak::BackendConnector::thread_starter(oWork); + printer::inst()->print_msg(L0, "Wait 30 sec until all backends are initialized"); + std::this_thread::sleep_for(std::chrono::seconds(30)); + + xmrstak::miner_work benchWork = xmrstak::miner_work("", work, sizeof(work), 0, false, 0); + printer::inst()->print_msg(L0, "Start a 60 second benchmark..."); + xmrstak::globalStates::inst().switch_work(benchWork, dat); uint64_t iStartStamp = get_timestamp_ms(); std::this_thread::sleep_for(std::chrono::seconds(60)); - - oWork = xmrstak::miner_work(); - xmrstak::pool_data dat; xmrstak::globalStates::inst().switch_work(oWork, dat); double fTotalHps = 0.0; @@ -734,9 +797,13 @@ void do_benchmark() double fHps = pvThreads->at(i)->iHashCount; fHps /= (pvThreads->at(i)->iTimestamp - iStartStamp) / 1000.0; - printer::inst()->print_msg(L0, "Thread %u: %.1f H/S", i, fHps); + auto bType = static_cast<xmrstak::iBackend::BackendType>(pvThreads->at(i)->backendType); + std::string name(xmrstak::iBackend::getName(bType)); + + printer::inst()->print_msg(L0, "Benchmark Thread %u %s: %.1f H/S", i,name.c_str(), fHps); fTotalHps += fHps; } - printer::inst()->print_msg(L0, "Total: %.1f H/S", fTotalHps); + printer::inst()->print_msg(L0, "Benchmark Total: %.1f H/S", fTotalHps); + return 0; } diff --git a/xmrstak/config.tpl b/xmrstak/config.tpl index 451ea7b..c95d142 100644 --- a/xmrstak/config.tpl +++ b/xmrstak/config.tpl @@ -1,28 +1,5 @@ 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. - * tls_fingerprint - Server's SHA256 fingerprint. If this string is non-empty then we will check the server's cert against it. - * pool_weight - Pool weight is a number telling the miner how important the pool is. Miner will mine mostly at the pool - * with the highest weight, unless the pool fails. Weight must be an integer larger than 0. - * - * We feature pools up to 1MH/s. For a more complete list see M5M400's pool list at www.moneropools.com - */ -"pool_list" : -[ -POOLCONF], - -/* - * currency to mine - * allowed values: 'monero' or 'aeon' - */ -"currency" : "CURRENCY", - -/* * Network timeouts. * Because of the way this client is written it doesn't need to constantly talk (keep-alive) to the server to make * sure it is there. We detect a buggy / overloaded server by the call timeout. The default values will be ok for diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp index c9d3a20..225fbe0 100644 --- a/xmrstak/jconf.cpp +++ b/xmrstak/jconf.cpp @@ -51,7 +51,7 @@ using namespace rapidjson; * This enum needs to match index in oConfigValues, otherwise we will get a runtime error */ enum configEnum { - aPoolList, bTlsSecureAlgo, sCurrency, iCallTimeout, iNetRetry, iGiveUpLimit, iVerboseLevel, bPrintMotd, iAutohashTime, + aPoolList, sCurrency, bTlsSecureAlgo, iCallTimeout, iNetRetry, iGiveUpLimit, iVerboseLevel, bPrintMotd, iAutohashTime, bFlushStdout, bDaemonMode, sOutputFile, iHttpdPort, sHttpLogin, sHttpPass, bPreferIpv4, bAesOverride, sUseSlowMem }; @@ -65,8 +65,8 @@ struct configVal { // kNullType means any type configVal oConfigValues[] = { { aPoolList, "pool_list", kArrayType }, - { bTlsSecureAlgo, "tls_secure_algo", kTrueType }, { sCurrency, "currency", kStringType }, + { bTlsSecureAlgo, "tls_secure_algo", kTrueType }, { iCallTimeout, "call_timeout", kNumberType }, { iNetRetry, "retry_time", kNumberType }, { iGiveUpLimit, "giveup_limit", kNumberType }, @@ -86,6 +86,28 @@ configVal oConfigValues[] = { constexpr size_t iConfigCnt = (sizeof(oConfigValues)/sizeof(oConfigValues[0])); +struct xmrstak_coin_algo +{ + const char* coin_name; + xmrstak_algo algo; + const char* default_pool; +}; + +xmrstak_coin_algo coin_algos[] = { + { "aeon", cryptonight_lite, "mine.aeon-pool.com:5555" }, + { "cryptonight", cryptonight, nullptr }, + { "cryptonight_lite", cryptonight_lite, nullptr }, + { "edollar", cryptonight, nullptr }, + { "electroneum", cryptonight, nullptr }, + { "graft", cryptonight, nullptr }, + { "intense", cryptonight, nullptr }, + { "karbo", cryptonight, nullptr }, + { "monero2", cryptonight_monero, "pool.usxmrpool.com:3333" }, + { "sumokoin", cryptonight_heavy, nullptr } +}; + +constexpr size_t coin_alogo_size = (sizeof(coin_algos)/sizeof(coin_algos[0])); + inline bool checkType(Type have, Type want) { if(want == have) @@ -103,6 +125,7 @@ inline bool checkType(Type have, Type want) struct jconf::opaque_private { Document jsonDoc; + Document jsonDocPools; const Value* configValues[iConfigCnt]; //Compile time constant opaque_private() @@ -168,45 +191,6 @@ bool jconf::TlsSecureAlgos() return prv->configValues[bTlsSecureAlgo]->GetBool(); } -const std::string jconf::GetCurrency() -{ - auto& currency = xmrstak::params::inst().currency; - if(currency.empty()) - currency = prv->configValues[sCurrency]->GetString(); - if( -#ifndef CONF_NO_MONERO - // if monero is disabled at compile time, enable error message if selected currency is `monero` - !xmrstak::strcmp_i(currency, "monero") -#else - true -#endif - && -#ifndef CONF_NO_AEON - // if aeon is disabled at compile time, enable error message if selected currency is `aeon` - !xmrstak::strcmp_i(currency, "aeon") -#else - true -#endif - ) - { - printer::inst()->print_msg(L0, "ERROR: Wrong currency selected - '%s'.", currency.c_str()); - win_exit(); - } - return currency; -} - -bool jconf::IsCurrencyMonero() -{ - if(xmrstak::strcmp_i(GetCurrency(), "monero")) - { - return true; - } - else - { - return false; - } -} - bool jconf::PreferIpv4() { return prv->configValues[bPreferIpv4]->GetBool(); @@ -312,18 +296,68 @@ jconf::slow_mem_cfg jconf::GetSlowMemSetting() return unknown_value; } -bool jconf::parse_config(const char* sFilename) +std::string jconf::GetMiningCoin() { - FILE * pFile; - char * buffer; - size_t flen; + if(xmrstak::params::inst().currency.length() > 0) + return xmrstak::params::inst().currency; + else + return prv->configValues[sCurrency]->GetString(); +} - if(!check_cpu_features()) +void jconf::GetAlgoList(std::string& list) +{ + list.reserve(256); + for(size_t i=0; i < coin_alogo_size; i++) { - printer::inst()->print_msg(L0, "CPU support of SSE2 is required."); + list += "\t- "; + list += coin_algos[i].coin_name; + list += "\n"; + } +} + +bool jconf::IsOnAlgoList(std::string& needle) +{ + std::transform(needle.begin(), needle.end(), needle.begin(), ::tolower); + + if(needle == "monero") + { + printer::inst()->print_msg(L0, "You entered Monero as coin name. Monero will hard-fork the PoW.\nThis means it will stop being compatible with other cryptonight coins.\n" + "Please use monero2 if you want to mine Monero, or name the coin that you want to mine."); return false; } + for(size_t i=0; i < coin_alogo_size; i++) + { + if(needle == coin_algos[i].coin_name) + return true; + } + return false; +} + +const char* jconf::GetDefaultPool(const char* needle) +{ + const char* default_example = "pool.example.com:3333"; + + for(size_t i=0; i < coin_alogo_size; i++) + { + if(strcmp(needle, coin_algos[i].coin_name) == 0) + { + if(coin_algos[i].default_pool != nullptr) + return coin_algos[i].default_pool; + else + return default_example; + } + } + + return default_example; +} + +bool jconf::parse_file(const char* sFilename, bool main_conf) +{ + FILE * pFile; + char * buffer; + size_t flen; + pFile = fopen(sFilename, "rb"); if (pFile == NULL) { @@ -372,46 +406,92 @@ bool jconf::parse_config(const char* sFilename) buffer[flen] = '}'; buffer[flen + 1] = '\0'; - prv->jsonDoc.Parse<kParseCommentsFlag|kParseTrailingCommasFlag>(buffer, flen+2); + Document& root = main_conf ? prv->jsonDoc : prv->jsonDocPools; + + root.Parse<kParseCommentsFlag|kParseTrailingCommasFlag>(buffer, flen+2); free(buffer); - if(prv->jsonDoc.HasParseError()) + if(root.HasParseError()) { - printer::inst()->print_msg(L0, "JSON config parse error(offset %llu): %s", - int_port(prv->jsonDoc.GetErrorOffset()), GetParseError_En(prv->jsonDoc.GetParseError())); + printer::inst()->print_msg(L0, "JSON config parse error in '%s' (offset %llu): %s", + sFilename, int_port(root.GetErrorOffset()), GetParseError_En(root.GetParseError())); return false; } - - if(!prv->jsonDoc.IsObject()) + if(!root.IsObject()) { //This should never happen as we created the root ourselves - printer::inst()->print_msg(L0, "Invalid config file. No root?\n"); + printer::inst()->print_msg(L0, "Invalid config file '%s'. No root?", sFilename); return false; } - for(size_t i = 0; i < iConfigCnt; i++) + if(main_conf) { - if(oConfigValues[i].iName != i) + for(size_t i = 2; i < iConfigCnt; i++) { - printer::inst()->print_msg(L0, "Code error. oConfigValues are not in order."); - return false; - } + if(oConfigValues[i].iName != i) + { + printer::inst()->print_msg(L0, "Code error. oConfigValues are not in order."); + return false; + } - prv->configValues[i] = GetObjectMember(prv->jsonDoc, oConfigValues[i].sName); + prv->configValues[i] = GetObjectMember(root, oConfigValues[i].sName); - if(prv->configValues[i] == nullptr) - { - printer::inst()->print_msg(L0, "Invalid config file. Missing value \"%s\".", oConfigValues[i].sName); - return false; - } + if(prv->configValues[i] == nullptr) + { + printer::inst()->print_msg(L0, "Invalid config file '%s'. Missing value \"%s\".", sFilename, oConfigValues[i].sName); + return false; + } - if(!checkType(prv->configValues[i]->GetType(), oConfigValues[i].iType)) + if(!checkType(prv->configValues[i]->GetType(), oConfigValues[i].iType)) + { + printer::inst()->print_msg(L0, "Invalid config file '%s'. Value \"%s\" has unexpected type.", sFilename, oConfigValues[i].sName); + return false; + } + } + } + else + { + for(size_t i = 0; i < 2; i++) { - printer::inst()->print_msg(L0, "Invalid config file. Value \"%s\" has unexpected type.", oConfigValues[i].sName); - return false; + if(oConfigValues[i].iName != i) + { + printer::inst()->print_msg(L0, "Code error. oConfigValues are not in order."); + return false; + } + + prv->configValues[i] = GetObjectMember(root, oConfigValues[i].sName); + + if(prv->configValues[i] == nullptr) + { + printer::inst()->print_msg(L0, "Invalid config file '%s'. Missing value \"%s\".", sFilename, oConfigValues[i].sName); + return false; + } + + if(!checkType(prv->configValues[i]->GetType(), oConfigValues[i].iType)) + { + printer::inst()->print_msg(L0, "Invalid config file '%s'. Value \"%s\" has unexpected type.", sFilename, oConfigValues[i].sName); + return false; + } } } + return true; +} + +bool jconf::parse_config(const char* sFilename, const char* sFilenamePools) +{ + if(!check_cpu_features()) + { + printer::inst()->print_msg(L0, "CPU support of SSE2 is required."); + return false; + } + + if(!parse_file(sFilename, true)) + return false; + + if(!parse_file(sFilenamePools, false)) + return false; + size_t pool_cnt = prv->configValues[aPoolList]->Size(); if(pool_cnt == 0) { @@ -529,5 +609,38 @@ bool jconf::parse_config(const char* sFilename) } } + std::string ctmp = GetMiningCoin(); + std::transform(ctmp.begin(), ctmp.end(), ctmp.begin(), ::tolower); + + if(ctmp.length() == 0) + { + printer::inst()->print_msg(L0, "You need to specify the coin that you want to mine."); + return false; + } + + for(size_t i=0; i < coin_alogo_size; i++) + { + if(ctmp == "monero") + { + printer::inst()->print_msg(L0, "You entered Monero as coin name. Monero will hard-fork the PoW.\nThis means it will stop being compatible with other cryptonight coins.\n" + "Please use monero2 if you want to mine Monero, or name the coin that you want to mine."); + return false; + } + + if(ctmp == coin_algos[i].coin_name) + { + mining_algo = coin_algos[i].algo; + break; + } + } + + if(mining_algo == invalid_algo) + { + std::string cl; + GetAlgoList(cl); + printer::inst()->print_msg(L0, "Unrecognised coin '%s', your options are:\n%s", ctmp.c_str(), cl.c_str()); + return false; + } + return true; } diff --git a/xmrstak/jconf.hpp b/xmrstak/jconf.hpp index 9a4e958..6874d37 100644 --- a/xmrstak/jconf.hpp +++ b/xmrstak/jconf.hpp @@ -1,12 +1,12 @@ #pragma once +#include "xmrstak/backend/cryptonight.hpp" #include "xmrstak/misc/environment.hpp" #include "params.hpp" #include <stdlib.h> #include <string> - class jconf { public: @@ -18,7 +18,7 @@ public: return env.pJconfConfig; }; - bool parse_config(const char* sFilename = xmrstak::params::inst().configFile.c_str()); + bool parse_config(const char* sFilename, const char* sFilenamePools); struct pool_cfg { const char* sPoolAddr; @@ -48,8 +48,13 @@ public: bool TlsSecureAlgos(); - const std::string GetCurrency(); - bool IsCurrencyMonero(); + inline xmrstak_algo GetMiningAlgo() { return mining_algo; } + + std::string GetMiningCoin(); + + static void GetAlgoList(std::string& list); + static bool IsOnAlgoList(std::string& needle); + static const char* GetDefaultPool(const char* needle); uint64_t GetVerboseLevel(); bool PrintMotd(); @@ -78,9 +83,12 @@ public: private: jconf(); + bool parse_file(const char* sFilename, bool main_conf); + bool check_cpu_features(); struct opaque_private; opaque_private* prv; bool bHaveAes; + xmrstak_algo mining_algo; }; diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp index c4ba26e..a620173 100644 --- a/xmrstak/misc/executor.cpp +++ b/xmrstak/misc/executor.cpp @@ -329,11 +329,14 @@ void executor::on_sock_ready(size_t pool_id) if(!pool->cmd_login()) { - if(!pool->have_sock_error()) + if(pool->have_call_error() && !pool->is_dev_pool()) { - log_socket_error(pool, pool->get_call_error()); - pool->disconnect(); + std::string str = "Login error: " + pool->get_call_error(); + log_socket_error(pool, std::move(str)); } + + if(!pool->have_sock_error()) + pool->disconnect(); } } @@ -406,14 +409,19 @@ void executor::on_pool_have_job(size_t pool_id, pool_job& oPoolJob) void executor::on_miner_result(size_t pool_id, job_result& oResult) { jpsock* pool = pick_pool_by_id(pool_id); - bool is_monero = jconf::inst()->IsCurrencyMonero(); + + const char* backend_name = xmrstak::iBackend::getName(pvThreads->at(oResult.iThreadId)->backendType); + uint64_t backend_hashcount, total_hashcount = 0; + + backend_hashcount = pvThreads->at(oResult.iThreadId)->iHashCount.load(std::memory_order_relaxed); + for(size_t i = 0; i < pvThreads->size(); i++) + total_hashcount += pvThreads->at(i)->iHashCount.load(std::memory_order_relaxed); if(pool->is_dev_pool()) { //Ignore errors silently if(pool->is_running() && pool->is_logged_in()) - pool->cmd_submit(oResult.sJobID, oResult.iNonce, oResult.bResult, pvThreads->at(oResult.iThreadId), is_monero); - + pool->cmd_submit(oResult.sJobID, oResult.iNonce, oResult.bResult, backend_name, backend_hashcount, total_hashcount, jconf::inst()->GetMiningAlgo()); return; } @@ -424,7 +432,7 @@ void executor::on_miner_result(size_t pool_id, job_result& oResult) } size_t t_start = get_timestamp_ms(); - bool bResult = pool->cmd_submit(oResult.sJobID, oResult.iNonce, oResult.bResult, pvThreads->at(oResult.iThreadId), is_monero); + bool bResult = pool->cmd_submit(oResult.sJobID, oResult.iNonce, oResult.bResult, backend_name, backend_hashcount, total_hashcount, jconf::inst()->GetMiningAlgo()); size_t t_len = get_timestamp_ms() - t_start; if(t_len > 0xFFFF) @@ -540,19 +548,38 @@ void executor::ex_main() 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()) + switch(jconf::inst()->GetMiningAlgo()) { + case cryptonight_heavy: 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:8888", "", "", "", 0.0, true, true, "", true); else - pools.emplace_front(0, "donate.xmr-stak.net:3333", "", "", "", 0.0, true, false, "", false); - } - else - { + pools.emplace_front(0, "donate.xmr-stak.net:5555", "", "", "", 0.0, true, false, "", true); + break; + + case cryptonight_monero: + if(dev_tls) + pools.emplace_front(0, "donate.xmr-stak.net:8800", "", "", "", 0.0, true, true, "", false); + else + pools.emplace_front(0, "donate.xmr-stak.net:5500", "", "", "", 0.0, true, false, "", false); + break; + + case cryptonight_lite: if(dev_tls) 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); + break; + + case cryptonight: + if(dev_tls) + 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); + break; + + default: + break; } ex_event ev; diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp index 9c413dc..95bcc9c 100644 --- a/xmrstak/net/jpsock.cpp +++ b/xmrstak/net/jpsock.cpp @@ -24,6 +24,7 @@ #include <stdarg.h> #include <assert.h> #include <algorithm> +#include <chrono> #include "jpsock.hpp" #include "socks.hpp" @@ -133,6 +134,7 @@ jpsock::~jpsock() std::string&& jpsock::get_call_error() { + call_error = false; return std::move(prv->oCallRsp.sCallErr); } @@ -189,11 +191,25 @@ bool jpsock::set_socket_error_strerr(const char* a, int res) void jpsock::jpsock_thread() { jpsock_thd_main(); + + if(!bHaveSocketError) + set_socket_error("Socket closed."); + executor::inst()->push_event(ex_event(std::move(sSocketError), quiet_close, pool_id)); - // If a call is wating, send an error to end it - bool bCallWaiting = false; std::unique_lock<std::mutex> mlock(call_mutex); + bool bWait = prv->oCallRsp.pCallData != nullptr; + + // If a call is waiting, wait a little bit before blowing it out of the water + if(bWait) + { + mlock.unlock(); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + mlock.lock(); + } + + // If the call is still there send an error to end it + bool bCallWaiting = false; if(prv->oCallRsp.pCallData != nullptr) { prv->oCallRsp.bHaveResponse = true; @@ -348,6 +364,7 @@ bool jpsock::process_line(char* line, size_t len) { prv->oCallRsp.pCallData = nullptr; prv->oCallRsp.sCallErr.assign(sError, iErrorLn); + call_error = true; } else prv->oCallRsp.pCallData->CopyFrom(*mt, prv->callAllocator); @@ -440,6 +457,7 @@ bool jpsock::connect(std::string& sConnectError) { ext_algo = ext_backend = ext_hashcount = ext_motd = false; bHaveSocketError = false; + call_error = false; sSocketError.clear(); iJobDiff = 0; connect_attempts++; @@ -596,7 +614,7 @@ bool jpsock::cmd_login() return true; } -bool jpsock::cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bResult, xmrstak::iBackend* bend, bool algo_full_cn) +bool jpsock::cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bResult, const char* backend_name, uint64_t backend_hashcount, uint64_t total_hashcount, xmrstak_algo algo) { char cmd_buffer[1024]; char sNonce[9]; @@ -604,16 +622,35 @@ bool jpsock::cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bRes /*Extensions*/ char sAlgo[64] = {0}; char sBackend[64] = {0}; - char sHashcount[64] = {0}; + char sHashcount[128] = {0}; if(ext_backend) - snprintf(sBackend, sizeof(sBackend), ",\"backend\":\"%s\"", xmrstak::iBackend::getName(bend->backendType)); + snprintf(sBackend, sizeof(sBackend), ",\"backend\":\"%s\"", backend_name); if(ext_hashcount) - snprintf(sHashcount, sizeof(sHashcount), ",\"hashcount\":%llu", int_port(bend->iHashCount.load(std::memory_order_relaxed))); + snprintf(sHashcount, sizeof(sHashcount), ",\"hashcount\":%llu,\"hashcount_total\":%llu", int_port(backend_hashcount), int_port(total_hashcount)); if(ext_algo) - snprintf(sAlgo, sizeof(sAlgo), ",\"algo\":\"%s\"", algo_full_cn ? "cryptonight" : "cryptonight-lite"); + { + const char* algo_name; + switch(algo) + { + case cryptonight: + algo_name = "cryptonight"; + break; + case cryptonight_lite: + algo_name = "cryptonight-lite"; + break; + case cryptonight_monero: + algo_name = "cryptonight-monero"; + break; + default: + algo_name = "unknown"; + break; + } + + snprintf(sAlgo, sizeof(sAlgo), ",\"algo\":\"%s\"", algo_name); + } bin2hex((unsigned char*)&iNonce, 4, sNonce); sNonce[8] = '\0'; diff --git a/xmrstak/net/jpsock.hpp b/xmrstak/net/jpsock.hpp index d9e5542..2ddeeee 100644 --- a/xmrstak/net/jpsock.hpp +++ b/xmrstak/net/jpsock.hpp @@ -2,6 +2,7 @@ #include "xmrstak/backend/iBackend.hpp" #include "msgstruct.hpp" +#include "xmrstak/jconf.hpp" #include <mutex> #include <atomic> @@ -34,7 +35,7 @@ public: void disconnect(bool quiet = false); bool cmd_login(); - bool cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bResult, xmrstak::iBackend* bend, bool algo_full_cn); + bool cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bResult, const char* backend_name, uint64_t backend_hashcount, uint64_t total_hashcount, xmrstak_algo algo); static bool hex2bin(const char* in, unsigned int len, unsigned char* out); static void bin2hex(const unsigned char* in, unsigned int len, char* out); @@ -62,6 +63,7 @@ public: bool get_pool_motd(std::string& strin); std::string&& get_call_error(); + bool have_call_error() { return call_error; } bool have_sock_error() { return bHaveSocketError; } inline static uint64_t t32_to_t64(uint32_t t) { return 0xFFFFFFFFFFFFFFFFULL / (0xFFFFFFFFULL / ((uint64_t)t)); } @@ -106,6 +108,7 @@ private: std::atomic<bool> bRunning; std::atomic<bool> bLoggedIn; std::atomic<bool> quiet_close; + std::atomic<bool> call_error; uint8_t* bJsonRecvMem; uint8_t* bJsonParseMem; diff --git a/xmrstak/net/socket.cpp b/xmrstak/net/socket.cpp index 89e9902..7c58a8e 100644 --- a/xmrstak/net/socket.cpp +++ b/xmrstak/net/socket.cpp @@ -48,6 +48,7 @@ bool plain_socket::set_hostname(const char* sAddr) char sAddrMb[256]; char *sTmp, *sPort; + sock_closed = false; size_t ln = strlen(sAddr); if (ln >= sizeof(sAddrMb)) return pCallback->set_socket_error("CONNECT error: Pool address overflow."); @@ -117,11 +118,16 @@ bool plain_socket::set_hostname(const char* sAddr) return pCallback->set_socket_error_strerr("CONNECT error: Socket creation failed "); } + int flag = 1; + /* If it fails, it fails, we won't loose too much sleep over it */ + setsockopt(hSocket, IPPROTO_TCP, TCP_NODELAY, (char *) &flag, sizeof(int)); + return true; } bool plain_socket::connect() { + sock_closed = false; int ret = ::connect(hSocket, pSockAddr->ai_addr, (int)pSockAddr->ai_addrlen); freeaddrinfo(pAddrRoot); @@ -135,6 +141,9 @@ bool plain_socket::connect() int plain_socket::recv(char* buf, unsigned int len) { + if(sock_closed) + return 0; + int ret = ::recv(hSocket, buf, len, 0); if(ret == 0) @@ -167,6 +176,7 @@ void plain_socket::close(bool free) { if(hSocket != INVALID_SOCKET) { + sock_closed = true; sock_close(hSocket); hSocket = INVALID_SOCKET; } @@ -211,12 +221,13 @@ void tls_socket::init_ctx() if(jconf::inst()->TlsSecureAlgos()) { - SSL_CTX_set_options(ctx, SSL_OP_NO_SSLv2 | SSL_OP_NO_SSLv3 | SSL_OP_NO_TLSv1 | SSL_OP_NO_COMPRESSION); + SSL_CTX_set_options(ctx, SSL_OP_NO_SSLv2 | SSL_OP_NO_SSLv3 | SSL_OP_NO_TLSv1); } } bool tls_socket::set_hostname(const char* sAddr) { + sock_closed = false; if(ctx == nullptr) { init_ctx(); @@ -233,6 +244,10 @@ bool tls_socket::set_hostname(const char* sAddr) return false; } + int flag = 1; + /* If it fails, it fails, we won't loose too much sleep over it */ + setsockopt(BIO_get_fd(bio, nullptr), IPPROTO_TCP, TCP_NODELAY, (char *) &flag, sizeof(int)); + if(BIO_set_conn_hostname(bio, sAddr) != 1) { print_error(); @@ -248,7 +263,7 @@ bool tls_socket::set_hostname(const char* sAddr) if(jconf::inst()->TlsSecureAlgos()) { - if(SSL_set_cipher_list(ssl, "HIGH:!aNULL:!kRSA:!PSK:!SRP:!MD5:!RC4:!SHA1") != 1) + if(SSL_set_cipher_list(ssl, "HIGH:!aNULL:!PSK:!SRP:!MD5:!RC4:!SHA1") != 1) { print_error(); return false; @@ -260,6 +275,7 @@ bool tls_socket::set_hostname(const char* sAddr) bool tls_socket::connect() { + sock_closed = false; if(BIO_do_connect(bio) != 1) { print_error(); @@ -340,6 +356,9 @@ bool tls_socket::connect() int tls_socket::recv(char* buf, unsigned int len) { + if(sock_closed) + return 0; + int ret = BIO_read(bio, buf, len); if(ret == 0) @@ -360,6 +379,7 @@ void tls_socket::close(bool free) if(bio == nullptr || ssl == nullptr) return; + sock_closed = true; if(!free) { sock_close(BIO_get_fd(bio, nullptr)); diff --git a/xmrstak/net/socket.hpp b/xmrstak/net/socket.hpp index 192a32c..b09142d 100644 --- a/xmrstak/net/socket.hpp +++ b/xmrstak/net/socket.hpp @@ -1,5 +1,6 @@ #pragma once +#include <atomic> #include "socks.hpp" class jpsock; @@ -12,6 +13,9 @@ public: virtual int recv(char* buf, unsigned int len) = 0; virtual bool send(const char* buf) = 0; virtual void close(bool free) = 0; + +protected: + std::atomic<bool> sock_closed; }; class plain_socket : public base_socket diff --git a/xmrstak/net/socks.hpp b/xmrstak/net/socks.hpp index 1d25d3a..86749e5 100644 --- a/xmrstak/net/socks.hpp +++ b/xmrstak/net/socks.hpp @@ -62,9 +62,8 @@ inline const char* sock_gai_strerror(int err, char* buf, size_t len) #include <unistd.h> /* Needed for close() */ #include <errno.h> #include <string.h> -#if defined(__FreeBSD__) #include <netinet/in.h> /* Needed for IPPROTO_TCP */ -#endif +#include <netinet/tcp.h> inline void sock_init() {} typedef int SOCKET; diff --git a/xmrstak/params.hpp b/xmrstak/params.hpp index bed3427..6928df5 100644 --- a/xmrstak/params.hpp +++ b/xmrstak/params.hpp @@ -40,6 +40,7 @@ struct params std::string currency; std::string configFile; + std::string configFilePools; std::string configFileAMD; std::string configFileNVIDIA; std::string configFileCPU; @@ -48,6 +49,9 @@ struct params std::string minerArg0; std::string minerArgs; + // block_version >= 0 enable benchmark + int benchmark_block_version = -1; + params() : binaryName("xmr-stak"), executablePrefix(""), @@ -55,6 +59,7 @@ struct params useNVIDIA(true), useCPU(true), configFile("config.txt"), + configFilePools("pools.txt"), configFileAMD("amd.txt"), configFileCPU("cpu.txt"), configFileNVIDIA("nvidia.txt") diff --git a/xmrstak/pools.tpl b/xmrstak/pools.tpl new file mode 100644 index 0000000..0b7084f --- /dev/null +++ b/xmrstak/pools.tpl @@ -0,0 +1,39 @@ +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. + * tls_fingerprint - Server's SHA256 fingerprint. If this string is non-empty then we will check the server's cert against it. + * pool_weight - Pool weight is a number telling the miner how important the pool is. Miner will mine mostly at the pool + * with the highest weight, unless the pool fails. Weight must be an integer larger than 0. + * + * We feature pools up to 1MH/s. For a more complete list see M5M400's pool list at www.moneropools.com + */ + +"pool_list" : +[ +POOLCONF], + +/* + * Currency to mine. Supported values: + * + * aeon + * cryptonight (try this if your coin is not listed) + * cryptonight_lite + * edollar + * electroneum + * graft + * intense + * karbo + * monero2 (use this for Monero's new PoW) + * sumokoin + * + */ + +"currency" : "CURRENCY", + +)===" + diff --git a/xmrstak/version.cpp b/xmrstak/version.cpp index 770834e..1444b30 100644 --- a/xmrstak/version.cpp +++ b/xmrstak/version.cpp @@ -32,13 +32,7 @@ #define OS_TYPE "unk" #endif -#if defined(CONF_NO_AEON) -#define COIN_TYPE "monero" -#elif defined(CONF_NO_MONERO) -#define COIN_TYPE "aeon" -#else -#define COIN_TYPE "aeon-monero" -#endif +#define COIN_TYPE "aeon-cryptonight-monero" #define XMRSTAK_PP_TOSTRING1(str) #str #define XMRSTAK_PP_TOSTRING(str) XMRSTAK_PP_TOSTRING1(str) |
