diff options
author | fireice-uk <fireice-uk@users.noreply.github.com> | 2018-04-02 21:14:14 +0100 |
---|---|---|
committer | GitHub <noreply@github.com> | 2018-04-02 21:14:14 +0100 |
commit | 6feffe2228282814abe716d241292ca2648a4e2a (patch) | |
tree | 277cc583c0f5df77e781dec0ba5019daf08e75d0 | |
parent | a036cd81592e3b3de804ba88bb8f94729ab60b7d (diff) | |
parent | 6880be70a6b29a1c850f4a4603bbc38d0ca6d117 (diff) | |
download | xmr-stak-6feffe2228282814abe716d241292ca2648a4e2a.zip xmr-stak-6feffe2228282814abe716d241292ca2648a4e2a.tar.gz |
Merge pull request #1279 from fireice-uk/dev
release 2.4.0
30 files changed, 290 insertions, 245 deletions
diff --git a/.travis.yml b/.travis.yml index de5b45f..f263e86 100644 --- a/.travis.yml +++ b/.travis.yml @@ -55,6 +55,22 @@ matrix: - CMAKE_C_COMPILER=gcc-6 - XMRSTAK_CMAKE_FLAGS="-DCUDA_ENABLE=OFF -DOpenCL_ENABLE=OFF" + # test with disabled HWLOC, MICROHTTPD, OpenSSL and no accelerators + - os: linux + compiler: gcc + addons: + apt: + sources: + - ubuntu-toolchain-r-test + packages: + - *default_packages + - gcc-6 + - g++-6 + env: + - CMAKE_CXX_COMPILER=g++-6 + - CMAKE_C_COMPILER=gcc-6 + - XMRSTAK_CMAKE_FLAGS="-DCUDA_ENABLE=OFF -DOpenCL_ENABLE=OFF -DHWLOC_ENABLE=OFF -DOpenSSL_ENABLE=OFF -DMICROHTTPD_ENABLE=OFF" + - os: linux compiler: gcc addons: @@ -1,6 +1,8 @@ ###### fireice-uk's and psychocrypt's # XMR-Stak - Monero/Aeon All-in-One Mining Software +**XMR-Stak is ready for the POW change of Monero-v7, Aeon-v7 and Sumukoin-v3** + XMR-Stak is a universal Stratum pool miner. This miner supports CPUs, AMD and NVIDIA gpus and can be used to mine the crypto currency Monero and Aeon. ## HTML reports @@ -8,14 +10,13 @@ XMR-Stak is a universal Stratum pool miner. This miner supports CPUs, AMD and NV ## Video setup guide on Windows -[<img src="https://gist.githubusercontent.com/fireice-uk/3621b179d56f57a8ead6303d8e415cf6/raw/4f2863d4072b78fdff649805e733203c9802daef/vidguidetmb.jpg">](https://www.youtube.com/watch?v=-8paGFwxyMU) +[<img src="https://gist.githubusercontent.com/fireice-uk/3621b179d56f57a8ead6303d8e415cf6/raw/f572faba67cc9418116f3c1dfd7783baf52182ce/vidguidetmb.jpg">](https://youtu.be/YNMa8NplWus) ###### Video by Crypto Sewer ## Overview * [Features](#features) * [Supported altcoins](#supported-altcoins) * [Download](#download) -* [Linux Portable Binary](doc/Linux_deployment.md) * [Usage](doc/usage.md) * [HowTo Compile](doc/compile.md) * [FAQ](doc/FAQ.md) @@ -57,7 +58,6 @@ Please note, this list is not complete, and is not an endorsement. ## Download You can find the latest releases and precompiled binaries on GitHub under [Releases](https://github.com/fireice-uk/xmr-stak/releases). -If you are running on Linux (especially Linux VMs), checkout [Linux Portable Binary](doc/Linux_deployment.md). ## Default Developer Donation @@ -72,7 +72,7 @@ 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 monero7` or `--currency aeon` + - start the miner with the [command line option](usage.md) `--currency monero7` or `--currency aeon7` - run `xmr-stak --help` to see all supported currencies and algorithms ## How can I mine Monero diff --git a/doc/Linux_deployment.md b/doc/Linux_deployment.md deleted file mode 100644 index 3219e8a..0000000 --- a/doc/Linux_deployment.md +++ /dev/null @@ -1,27 +0,0 @@ -# Deploying portable **XMR-Stak** on Linux systems - -**This is an experimental feature** we reserve the right to remove the binary if we get too many issues. - -XMR-Stak releases include a pre-built portable version. If you are simply using it to avoid having to compile the application, you can simply download **xmr-stak-portbin-linux.tar.gz** from our [latest releases](https://github.com/fireice-uk/xmr-stak/releases/latest). Open up command line, and use the following commands: - -``` -tar xzf xmr-stak-portbin-linux.tar.gz -./xmr-stak.sh -``` - -Configuration and tuning files will be generated automatically from your answers. - -For automatic deployments, please use the steps above to obtain config.txt and use the following script: - -``` -#!/bin/bash -sudo apt install curl -curl -O `curl -s https://api.github.com/repos/fireice-uk/xmr-stak/releases/latest | grep -o 'browser_download_url.*xmr-stak-portbin-linux.tar.gz' | sed 's/.*"//'` -curl -O http://path.to/your/config.txt -tar xzf xmr-stak-portbin-linux.tar.gz -./xmr-stak.sh -``` - -XMR-Stak will auto-configure and go to work. You don't even need Docker! - - diff --git a/doc/compile.md b/doc/compile.md index 984c013..2f9dace 100644 --- a/doc/compile.md +++ b/doc/compile.md @@ -35,8 +35,8 @@ After the configuration you need to compile the miner, follow the guide for your ## Generic Build Options - `CMAKE_INSTALL_PREFIX` install miner to the home folder - - `cmake .. -DCMAKE_INSTALL_PREFIX=$HOME/xmr-stak-cpu` - - you can find the binary and the `config.txt` file after `make install` in `$HOME/xmr-stak-cpu/bin` + - `cmake .. -DCMAKE_INSTALL_PREFIX=$HOME/xmr-stak` + - you can find the binary and the `config.txt` file after `make install` in `$HOME/xmr-stak/bin` - `CMAKE_LINK_STATIC` link libgcc and libstdc++ libraries static (default OFF) - disable with `cmake .. -DCMAKE_LINK_STATIC=ON` - if you use static compile to run the miner on another system set `-DXMR-STAK_COMPILE=generic` @@ -55,7 +55,7 @@ After the configuration you need to compile the miner, follow the guide for your - `CPU_ENABLE` allow to disable/enable the CPU backend of the miner - `HWLOC_ENABLE` allow to disable/enable the dependency *hwloc* - - the config suggestion is not optimal if option is disabled: `cmake . -DHWLOC_ENABLE=OFF` + - the config suggestion is not optimal if option is disabled: `cmake .. -DHWLOC_ENABLE=OFF` - disabling can be reduce the miner performance ## AMD Build Options diff --git a/doc/compile_Linux.md b/doc/compile_Linux.md index 729f4d2..3eb4960 100644 --- a/doc/compile_Linux.md +++ b/doc/compile_Linux.md @@ -97,6 +97,13 @@ - g++ version 5.1 or higher is required for full C++11 support. If you want to compile the binary without installing libraries / compiler or just compile binary for some other distribution, please check the [build_xmr-stak_docker.sh script](scripts/build_xmr-stak_docker/build_xmr-stak_docker.sh). +- Some newer gcc versions are not supported by CUDA (e.g. Ubuntu 17.10). It will require installing gcc 5 but you can avoid changing defaults. + +In that case you can force CUDA to use an older compiler in the following way: +``` +cmake -DCUDA_HOST_COMPILER=/usr/bin/gcc-5 .. +``` + ### To do a generic and static build for a system without gcc 5.1+ ``` cmake -DCMAKE_LINK_STATIC=ON -DXMR-STAK_COMPILE=generic . diff --git a/doc/usage.md b/doc/usage.md index 1f1fb09..886c1b3 100644 --- a/doc/usage.md +++ b/doc/usage.md @@ -16,6 +16,7 @@ The number of files depends on the available backends. `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)). +Note: If the pool is ignoring the option `rig_id` in `pools.txt` to name your worker please check the pool documentation how a worker name can be set. ## Usage on Windows 1) Double click the `xmr-stak.exe` file diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index 8d0fd32..2973db4 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -16,6 +16,7 @@ #include "xmrstak/backend/cryptonight.hpp" #include "xmrstak/jconf.hpp" #include "xmrstak/picosha2/picosha2.hpp" +#include "xmrstak/params.hpp" #include <stdio.h> #include <string.h> @@ -393,9 +394,10 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_ std::string cache_file = get_home() + "/.openclcache/" + hash_hex_str + ".openclbin"; std::ifstream clBinFile(cache_file, std::ofstream::in | std::ofstream::binary); - if(!clBinFile.good()) + if(xmrstak::params::inst().AMDCache == false || !clBinFile.good()) { - printer::inst()->print_msg(L1,"OpenCL device %u - Precompiled code %s not found. Compiling ...",ctx->deviceIdx, cache_file.c_str()); + if(xmrstak::params::inst().AMDCache) + printer::inst()->print_msg(L1,"OpenCL device %u - Precompiled code %s not found. Compiling ...",ctx->deviceIdx, cache_file.c_str()); ctx->Program = clCreateProgramWithSource(opencl_ctx, 1, (const char**)&source_code, NULL, &ret); if(ret != CL_SUCCESS) { @@ -467,29 +469,31 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_ std::vector<char*> all_programs(num_devices); std::vector<std::vector<char>> program_storage; - int p_id = 0; - size_t mem_size = 0; - // create memory structure to query all OpenCL program binaries - for(auto & p : all_programs) + if(xmrstak::params::inst().AMDCache) { - program_storage.emplace_back(std::vector<char>(binary_sizes[p_id])); - all_programs[p_id] = program_storage[p_id].data(); - mem_size += binary_sizes[p_id]; - p_id++; - } + int p_id = 0; + size_t mem_size = 0; + // create memory structure to query all OpenCL program binaries + for(auto & p : all_programs) + { + program_storage.emplace_back(std::vector<char>(binary_sizes[p_id])); + all_programs[p_id] = program_storage[p_id].data(); + mem_size += binary_sizes[p_id]; + p_id++; + } - if( ret = clGetProgramInfo(ctx->Program, CL_PROGRAM_BINARIES, num_devices * sizeof(char*), all_programs.data(),NULL) != CL_SUCCESS) - { - printer::inst()->print_msg(L1,"Error %s when calling clGetProgramInfo.", err_to_str(ret)); - return ERR_OCL_API; - } + if((ret = clGetProgramInfo(ctx->Program, CL_PROGRAM_BINARIES, num_devices * sizeof(char*), all_programs.data(),NULL)) != CL_SUCCESS) + { + printer::inst()->print_msg(L1,"Error %s when calling clGetProgramInfo.", err_to_str(ret)); + return ERR_OCL_API; + } - std::ofstream file_stream; - std::cout<<get_home() + "/.openclcache/" + hash_hex_str + ".openclbin"<<std::endl; - file_stream.open(cache_file, std::ofstream::out | std::ofstream::binary); - file_stream.write(all_programs[dev_id], binary_sizes[dev_id]); - file_stream.close(); - printer::inst()->print_msg(L1, "OpenCL device %u - Precompiled code stored in file %s",ctx->deviceIdx, cache_file.c_str()); + std::ofstream file_stream; + file_stream.open(cache_file, std::ofstream::out | std::ofstream::binary); + file_stream.write(all_programs[dev_id], binary_sizes[dev_id]); + file_stream.close(); + printer::inst()->print_msg(L1, "OpenCL device %u - Precompiled code stored in file %s",ctx->deviceIdx, cache_file.c_str()); + } } else { @@ -936,7 +940,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar /// @todo only activate if currency is monero int cn_kernel_offset = 0; - if(miner_algo == cryptonight_monero && version >= 7) + if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) { cn_kernel_offset = 6; } @@ -962,7 +966,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar return(ERR_OCL_API); } - if(miner_algo == cryptonight_monero && version >= 7) + if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon ) { // Input if ((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS) @@ -1130,7 +1134,7 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo, size_t tmpNonce = ctx->Nonce; /// @todo only activate if currency is monero int cn_kernel_offset = 0; - if(miner_algo == cryptonight_monero && version >= 7) + if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon) { cn_kernel_offset = 6; } diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl index 7a36357..5d4e66c 100644 --- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl +++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl @@ -433,10 +433,7 @@ 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)]; -} +#define mix_and_propagate(xin) (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 diff --git a/xmrstak/backend/amd/jconf.cpp b/xmrstak/backend/amd/jconf.cpp index 93ba709..0f39ff2 100644 --- a/xmrstak/backend/amd/jconf.cpp +++ b/xmrstak/backend/amd/jconf.cpp @@ -234,15 +234,15 @@ bool jconf::parse_config(const char* sFilename) if(prv->jsonDoc.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(prv->jsonDoc.GetErrorOffset()), GetParseError_En(prv->jsonDoc.GetParseError())); return false; } if(!prv->jsonDoc.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; } @@ -258,13 +258,13 @@ bool jconf::parse_config(const char* sFilename) if(prv->configValues[i] == nullptr) { - printer::inst()->print_msg(L0, "Invalid config file. Missing value \"%s\".", oConfigValues[i].sName); + 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. Value \"%s\" has unexpected type.", oConfigValues[i].sName); + printer::inst()->print_msg(L0, "Invalid config file '%s'. Value \"%s\" has unexpected type.", sFilename, oConfigValues[i].sName); return false; } } diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp index 46a04d5..cab5ad9 100644 --- a/xmrstak/backend/amd/minethd.cpp +++ b/xmrstak/backend/amd/minethd.cpp @@ -191,15 +191,10 @@ void minethd::work_main() uint64_t iCount = 0; cryptonight_ctx* cpu_ctx; cpu_ctx = cpu::minethd::minethd_alloc_ctx(); - 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); + + // start with root algorithm and switch later if fork version is reached + auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot(); + cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); globalStates::inst().iConsumeCnt++; @@ -218,17 +213,18 @@ 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) + continue; + } + + uint8_t new_version = oWork.getVersion(); + if(new_version != version) + { + if(new_version >= ::jconf::inst()->GetMiningForkVersion()) { - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy); + miner_algo = ::jconf::inst()->GetMiningAlgo(); + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); } version = new_version; - continue; } uint32_t h_per_round = pGpuCtx->rawIntensity; @@ -236,7 +232,7 @@ void minethd::work_main() assert(sizeof(job_result::sJobID) == sizeof(pool_job::sJobID)); uint64_t target = oWork.iTarget; - /// \todo add monero hard for version + XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target, miner_algo, version); if(oWork.bNiceHash) @@ -280,16 +276,6 @@ 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/autoAdjust.hpp b/xmrstak/backend/cpu/autoAdjust.hpp index db805ec..969d478 100644 --- a/xmrstak/backend/cpu/autoAdjust.hpp +++ b/xmrstak/backend/cpu/autoAdjust.hpp @@ -35,19 +35,9 @@ public: bool printConfig() { - size_t hashMemSizeKB; - size_t halfHashMemSizeKB; - if(::jconf::inst()->IsCurrencyMonero()) - { - hashMemSizeKB = MONERO_MEMORY / 1024u; - halfHashMemSizeKB = hashMemSizeKB / 2u; - } - else - { - hashMemSizeKB = AEON_MEMORY / 1024u; - halfHashMemSizeKB = hashMemSizeKB / 2u; - } + const size_t hashMemSizeKB = cn_select_memory(::jconf::inst()->GetMiningAlgo()) / 1024u; + const size_t halfHashMemSizeKB = hashMemSizeKB / 2u; configEditor configTpl{}; diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h index 85373e8..5203de8 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h +++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h @@ -444,7 +444,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c constexpr size_t ITERATIONS = cn_select_iter<ALGO>(); constexpr size_t MEM = cn_select_memory<ALGO>(); - if(ALGO == cryptonight_monero && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43) { memset(output, 0, 32); return; @@ -453,7 +453,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c keccak((const uint8_t *)input, len, ctx0->hash_state, 200); uint64_t monero_const; - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) { 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); @@ -482,7 +482,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c else cx = _mm_aesenc_si128(cx, _mm_set_epi64x(ah0, al0)); - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) 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)); @@ -506,7 +506,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c _mm_prefetch((const char*)&l0[al0 & MASK], _MM_HINT_T0); ah0 += lo; - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) ((uint64_t*)&l0[idx0 & MASK])[1] = ah0 ^ monero_const; else ((uint64_t*)&l0[idx0 & MASK])[1] = ah0; @@ -544,7 +544,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto constexpr size_t ITERATIONS = cn_select_iter<ALGO>(); constexpr size_t MEM = cn_select_memory<ALGO>(); - if(ALGO == cryptonight_monero && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43) { memset(output, 0, 64); return; @@ -554,7 +554,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto keccak((const uint8_t *)input+len, len, ctx[1]->hash_state, 200); uint64_t monero_const_0, monero_const_1; - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) { 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); @@ -592,7 +592,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto else cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh0, axl0)); - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) 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)); @@ -610,7 +610,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto else cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh1, axl1)); - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) 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)); @@ -631,7 +631,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto axh0 += lo; ((uint64_t*)&l0[idx0 & MASK])[0] = axl0; - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) ((uint64_t*)&l0[idx0 & MASK])[1] = axh0 ^ monero_const_0; else ((uint64_t*)&l0[idx0 & MASK])[1] = axh0; @@ -662,7 +662,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto axh1 += lo; ((uint64_t*)&l1[idx1 & MASK])[0] = axl1; - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) ((uint64_t*)&l1[idx1 & MASK])[1] = axh1 ^ monero_const_1; else ((uint64_t*)&l1[idx1 & MASK])[1] = axh1; @@ -709,7 +709,7 @@ 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); \ - if(ALGO == cryptonight_monero) \ + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) \ cryptonight_monero_tweak((uint64_t*)ptr, b); \ else \ _mm_store_si128(ptr, b);\ @@ -724,7 +724,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto #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)); \ - if(ALGO == cryptonight_monero) \ + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) \ _mm_store_si128(ptr, _mm_xor_si128(a, mc)); \ else \ _mm_store_si128(ptr, a);\ @@ -751,7 +751,7 @@ void cryptonight_triple_hash(const void* input, size_t len, void* output, crypto constexpr size_t ITERATIONS = cn_select_iter<ALGO>(); constexpr size_t MEM = cn_select_memory<ALGO>(); - if(ALGO == cryptonight_monero && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43) { memset(output, 0, 32 * 3); return; @@ -845,7 +845,7 @@ void cryptonight_quad_hash(const void* input, size_t len, void* output, cryptoni constexpr size_t ITERATIONS = cn_select_iter<ALGO>(); constexpr size_t MEM = cn_select_memory<ALGO>(); - if(ALGO == cryptonight_monero && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43) { memset(output, 0, 32 * 4); return; @@ -954,7 +954,7 @@ void cryptonight_penta_hash(const void* input, size_t len, void* output, crypton constexpr size_t ITERATIONS = cn_select_iter<ALGO>(); constexpr size_t MEM = cn_select_memory<ALGO>(); - if(ALGO == cryptonight_monero && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43) { memset(output, 0, 32 * 5); return; diff --git a/xmrstak/backend/cpu/jconf.cpp b/xmrstak/backend/cpu/jconf.cpp index 6e709bd..399dd16 100644 --- a/xmrstak/backend/cpu/jconf.cpp +++ b/xmrstak/backend/cpu/jconf.cpp @@ -211,14 +211,14 @@ bool jconf::parse_config(const char* sFilename) if(prv->jsonDoc.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(prv->jsonDoc.GetErrorOffset()), GetParseError_En(prv->jsonDoc.GetParseError())); return false; } if(!prv->jsonDoc.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; } @@ -234,13 +234,13 @@ bool jconf::parse_config(const char* sFilename) if(prv->configValues[i] == nullptr) { - printer::inst()->print_msg(L0, "Invalid config file. Missing value \"%s\".", oConfigValues[i].sName); + 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. Value \"%s\" has unexpected type.", oConfigValues[i].sName); + printer::inst()->print_msg(L0, "Invalid config file '%s'. Value \"%s\" has unexpected type.", sFilename, oConfigValues[i].sName); return false; } } diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp index e263aca..5b56f85 100644 --- a/xmrstak/backend/cpu/minethd.cpp +++ b/xmrstak/backend/cpu/minethd.cpp @@ -282,6 +282,9 @@ bool minethd::self_test() else if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero) { } + else if(::jconf::inst()->GetMiningAlgo() == cryptonight_aeon) + { + } for (int i = 0; i < MAX_N; i++) cryptonight_free_ctx(ctx[i]); @@ -366,6 +369,9 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr case cryptonight_heavy: algv = 3; break; + case cryptonight_aeon: + algv = 4; + break; default: algv = 2; break; @@ -387,7 +393,11 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr cryptonight_hash<cryptonight_heavy, false, false>, cryptonight_hash<cryptonight_heavy, true, false>, cryptonight_hash<cryptonight_heavy, false, true>, - cryptonight_hash<cryptonight_heavy, true, true> + cryptonight_hash<cryptonight_heavy, true, true>, + cryptonight_hash<cryptonight_aeon, false, false>, + cryptonight_hash<cryptonight_aeon, true, false>, + cryptonight_hash<cryptonight_aeon, false, true>, + cryptonight_hash<cryptonight_aeon, true, true> }; std::bitset<2> digit; @@ -407,14 +417,15 @@ void minethd::work_main() lck.release(); std::this_thread::yield(); - cn_hash_fun hash_fun; cryptonight_ctx* ctx; uint64_t iCount = 0; uint64_t* piHashVal; uint32_t* piNonce; job_result result; - hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo()); + // start with root algorithm and switch later if fork version is reached + auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot(); + cn_hash_fun hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo); ctx = minethd_alloc_ctx(); piHashVal = (uint64_t*)(result.bResult + 24); @@ -422,6 +433,8 @@ void minethd::work_main() globalStates::inst().inst().iConsumeCnt++; result.iThreadId = iThreadNo; + uint8_t version = 0; + while (bQuit == 0) { if (oWork.bStall) @@ -447,20 +460,15 @@ void minethd::work_main() if(oWork.bNiceHash) result.iNonce = *piNonce; - if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero) + uint8_t new_version = oWork.getVersion(); + if(new_version != version) { - 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); + if(new_version >= ::jconf::inst()->GetMiningForkVersion()) + { + miner_algo = ::jconf::inst()->GetMiningAlgo(); + hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo); + } + version = new_version; } while(globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo) @@ -512,6 +520,12 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, case cryptonight_monero: algv = 0; break; + case cryptonight_heavy: + algv = 3; + break; + case cryptonight_aeon: + algv = 4; + break; default: algv = 2; break; @@ -534,6 +548,7 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, 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>, @@ -550,6 +565,7 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, 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>, @@ -565,7 +581,41 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, cryptonight_penta_hash<cryptonight, false, false>, cryptonight_penta_hash<cryptonight, true, false>, cryptonight_penta_hash<cryptonight, false, true>, - cryptonight_penta_hash<cryptonight, true, true> + cryptonight_penta_hash<cryptonight, true, true>, + + cryptonight_double_hash<cryptonight_heavy, false, false>, + cryptonight_double_hash<cryptonight_heavy, true, false>, + cryptonight_double_hash<cryptonight_heavy, false, true>, + cryptonight_double_hash<cryptonight_heavy, true, true>, + cryptonight_triple_hash<cryptonight_heavy, false, false>, + cryptonight_triple_hash<cryptonight_heavy, true, false>, + cryptonight_triple_hash<cryptonight_heavy, false, true>, + cryptonight_triple_hash<cryptonight_heavy, true, true>, + cryptonight_quad_hash<cryptonight_heavy, false, false>, + cryptonight_quad_hash<cryptonight_heavy, true, false>, + cryptonight_quad_hash<cryptonight_heavy, false, true>, + cryptonight_quad_hash<cryptonight_heavy, true, true>, + cryptonight_penta_hash<cryptonight_heavy, false, false>, + cryptonight_penta_hash<cryptonight_heavy, true, false>, + cryptonight_penta_hash<cryptonight_heavy, false, true>, + cryptonight_penta_hash<cryptonight_heavy, true, true>, + + cryptonight_double_hash<cryptonight_aeon, false, false>, + cryptonight_double_hash<cryptonight_aeon, true, false>, + cryptonight_double_hash<cryptonight_aeon, false, true>, + cryptonight_double_hash<cryptonight_aeon, true, true>, + cryptonight_triple_hash<cryptonight_aeon, false, false>, + cryptonight_triple_hash<cryptonight_aeon, true, false>, + cryptonight_triple_hash<cryptonight_aeon, false, true>, + cryptonight_triple_hash<cryptonight_aeon, true, true>, + cryptonight_quad_hash<cryptonight_aeon, false, false>, + cryptonight_quad_hash<cryptonight_aeon, true, false>, + cryptonight_quad_hash<cryptonight_aeon, false, true>, + cryptonight_quad_hash<cryptonight_aeon, true, true>, + cryptonight_penta_hash<cryptonight_aeon, false, false>, + cryptonight_penta_hash<cryptonight_aeon, true, false>, + cryptonight_penta_hash<cryptonight_aeon, false, true>, + cryptonight_penta_hash<cryptonight_aeon, true, true> }; std::bitset<2> digit; @@ -577,22 +627,22 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, void minethd::double_work_main() { - multiway_work_main<2>(func_multi_selector(2, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo())); + multiway_work_main<2>(); } void minethd::triple_work_main() { - multiway_work_main<3>(func_multi_selector(3, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo())); + multiway_work_main<3>(); } void minethd::quad_work_main() { - multiway_work_main<4>(func_multi_selector(4, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo())); + multiway_work_main<4>(); } void minethd::penta_work_main() { - multiway_work_main<5>(func_multi_selector(5, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, ::jconf::inst()->GetMiningAlgo())); + multiway_work_main<5>(); } template<size_t N> @@ -607,7 +657,7 @@ void minethd::prep_multiway_work(uint8_t *bWorkBlob, uint32_t **piNonce) } template<size_t N> -void minethd::multiway_work_main(cn_hash_fun_multi hash_fun_multi) +void minethd::multiway_work_main() { if(affinity >= 0) //-1 means no affinity bindMemoryToNUMANode(affinity); @@ -638,6 +688,11 @@ void minethd::multiway_work_main(cn_hash_fun_multi hash_fun_multi) globalStates::inst().iConsumeCnt++; + // start with root algorithm and switch later if fork version is reached + auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot(); + cn_hash_fun_multi hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo); + uint8_t version = 0; + while (bQuit == 0) { if (oWork.bStall) @@ -662,20 +717,15 @@ void minethd::multiway_work_main(cn_hash_fun_multi hash_fun_multi) if(oWork.bNiceHash) iNonce = *piNonce[0]; - if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero) + uint8_t new_version = oWork.getVersion(); + if(new_version != version) { - 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); + if(new_version >= ::jconf::inst()->GetMiningForkVersion()) + { + miner_algo = ::jconf::inst()->GetMiningAlgo(); + hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo); + } + version = new_version; } while (globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo) diff --git a/xmrstak/backend/cpu/minethd.hpp b/xmrstak/backend/cpu/minethd.hpp index ef1bbd2..cd78343 100644 --- a/xmrstak/backend/cpu/minethd.hpp +++ b/xmrstak/backend/cpu/minethd.hpp @@ -36,7 +36,7 @@ private: minethd(miner_work& pWork, size_t iNo, int iMultiway, bool no_prefetch, int64_t affinity); template<size_t N> - void multiway_work_main(cn_hash_fun_multi hash_fun_multi); + void multiway_work_main(); template<size_t N> void prep_multiway_work(uint8_t *bWorkBlob, uint32_t **piNonce); diff --git a/xmrstak/backend/cryptonight.hpp b/xmrstak/backend/cryptonight.hpp index fe10a9f..8a8e259 100644 --- a/xmrstak/backend/cryptonight.hpp +++ b/xmrstak/backend/cryptonight.hpp @@ -9,7 +9,8 @@ enum xmrstak_algo cryptonight = 1, cryptonight_lite = 2, cryptonight_monero = 3, - cryptonight_heavy = 4 + cryptonight_heavy = 4, + cryptonight_aeon = 5 }; // define aeon settings @@ -40,17 +41,20 @@ inline constexpr size_t cn_select_memory<cryptonight_monero>() { return CRYPTONI template<> inline constexpr size_t cn_select_memory<cryptonight_heavy>() { return CRYPTONIGHT_HEAVY_MEMORY; } +template<> +inline constexpr size_t cn_select_memory<cryptonight_aeon>() { return CRYPTONIGHT_LITE_MEMORY; } + inline size_t cn_select_memory(xmrstak_algo algo) { switch(algo) { + case cryptonight_monero: case cryptonight: - return CRYPTONIGHT_MEMORY; + return CRYPTONIGHT_MEMORY; + case cryptonight_aeon: case cryptonight_lite: return CRYPTONIGHT_LITE_MEMORY; - case cryptonight_monero: - return CRYPTONIGHT_MEMORY; case cryptonight_heavy: return CRYPTONIGHT_HEAVY_MEMORY; default: @@ -73,16 +77,19 @@ inline constexpr uint32_t cn_select_mask<cryptonight_monero>() { return CRYPTONI template<> inline constexpr uint32_t cn_select_mask<cryptonight_heavy>() { return CRYPTONIGHT_HEAVY_MASK; } +template<> +inline constexpr uint32_t cn_select_mask<cryptonight_aeon>() { return CRYPTONIGHT_LITE_MASK; } + inline size_t cn_select_mask(xmrstak_algo algo) { switch(algo) { + case cryptonight_monero: case cryptonight: return CRYPTONIGHT_MASK; + case cryptonight_aeon: case cryptonight_lite: return CRYPTONIGHT_LITE_MASK; - case cryptonight_monero: - return CRYPTONIGHT_MASK; case cryptonight_heavy: return CRYPTONIGHT_HEAVY_MASK; default: @@ -105,16 +112,19 @@ inline constexpr uint32_t cn_select_iter<cryptonight_monero>() { return CRYPTONI template<> inline constexpr uint32_t cn_select_iter<cryptonight_heavy>() { return CRYPTONIGHT_HEAVY_ITER; } +template<> +inline constexpr uint32_t cn_select_iter<cryptonight_aeon>() { return CRYPTONIGHT_LITE_ITER; } + inline size_t cn_select_iter(xmrstak_algo algo) { switch(algo) { + case cryptonight_monero: case cryptonight: return CRYPTONIGHT_ITER; + case cryptonight_aeon: case cryptonight_lite: return CRYPTONIGHT_LITE_ITER; - case cryptonight_monero: - return CRYPTONIGHT_ITER; case cryptonight_heavy: return CRYPTONIGHT_HEAVY_ITER; default: diff --git a/xmrstak/backend/nvidia/jconf.cpp b/xmrstak/backend/nvidia/jconf.cpp index 46c5726..c9d4f19 100644 --- a/xmrstak/backend/nvidia/jconf.cpp +++ b/xmrstak/backend/nvidia/jconf.cpp @@ -235,15 +235,15 @@ bool jconf::parse_config(const char* sFilename) if(prv->jsonDoc.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(prv->jsonDoc.GetErrorOffset()), GetParseError_En(prv->jsonDoc.GetParseError())); return false; } if(!prv->jsonDoc.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; } @@ -259,13 +259,13 @@ bool jconf::parse_config(const char* sFilename) if(prv->configValues[i] == nullptr) { - printer::inst()->print_msg(L0, "Invalid config file. Missing value \"%s\".", oConfigValues[i].sName); + 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. Value \"%s\" has unexpected type.", oConfigValues[i].sName); + printer::inst()->print_msg(L0, "Invalid config file '%s'. Value \"%s\" has unexpected type.", sFilename, oConfigValues[i].sName); return false; } } @@ -274,4 +274,4 @@ bool jconf::parse_config(const char* sFilename) } } // namespace nvidia -} // namespace xmrstak
\ No newline at end of file +} // namespace xmrstak diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp index 153e4e3..804c06a 100644 --- a/xmrstak/backend/nvidia/minethd.cpp +++ b/xmrstak/backend/nvidia/minethd.cpp @@ -237,15 +237,11 @@ void minethd::work_main() uint64_t iCount = 0; cryptonight_ctx* cpu_ctx; cpu_ctx = cpu::minethd::minethd_alloc_ctx(); - 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); + + // start with root algorithm and switch later if fork version is reached + auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot(); + cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); + uint32_t iNonce; globalStates::inst().iConsumeCnt++; @@ -265,17 +261,17 @@ 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) + continue; + } + uint8_t new_version = oWork.getVersion(); + if(new_version != version) + { + if(new_version >= ::jconf::inst()->GetMiningForkVersion()) { - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy); + miner_algo = ::jconf::inst()->GetMiningAlgo(); + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); } version = new_version; - continue; } cryptonight_extra_cpu_set_data(&ctx, oWork.bWorkBlob, oWork.iWorkSize); @@ -299,11 +295,11 @@ void minethd::work_main() uint32_t foundNonce[10]; uint32_t foundCount; - cryptonight_extra_cpu_prepare(&ctx, iNonce, miner_algo, version); + cryptonight_extra_cpu_prepare(&ctx, iNonce, miner_algo); - cryptonight_core_cpu_hash(&ctx, miner_algo, iNonce, version); + cryptonight_core_cpu_hash(&ctx, miner_algo, iNonce); - cryptonight_extra_cpu_final(&ctx, iNonce, oWork.iTarget, &foundCount, foundNonce, miner_algo, version); + cryptonight_extra_cpu_final(&ctx, iNonce, oWork.iTarget, &foundCount, foundNonce, miner_algo); for(size_t i = 0; i < foundCount; i++) { @@ -334,19 +330,8 @@ 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; } } } // namespace xmrstak - } //namespace nvidia diff --git a/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp b/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp index 29a3523..c1e31b9 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp +++ b/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp @@ -45,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, 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_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce, xmrstak_algo miner_algo); +void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce,xmrstak_algo miner_algo); } -void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t startNonce, uint8_t version); +void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t startNonce); diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu index ede578f..7aa44e8 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu @@ -231,7 +231,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti uint32_t t1[2], t2[2], res; uint32_t tweak1_2[2]; - if (ALGO == cryptonight_monero) + if (ALGO == cryptonight_monero || ALGO == cryptonight_aeon) { uint32_t * state = d_ctx_state + thread * 50; tweak1_2[0] = (d_input[8] >> 24) | (d_input[9] << 8); @@ -275,7 +275,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti t1[0] = shuffle<4>(sPtr,sub, d[x], 0); const uint32_t z = d[0] ^ d[1]; - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) { const uint32_t table = 0x75310U; const uint32_t index = ((z >> 26) & 12) | ((z >> 23) & 2); @@ -304,7 +304,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti res = *( (uint64_t *) t2 ) >> ( sub & 1 ? 32 : 0 ); - if(ALGO == cryptonight_monero) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) { const uint32_t tweaked_res = tweak1_2[sub & 1] ^ res; const uint32_t long_state_update = sub2 ? tweaked_res : res; @@ -475,24 +475,16 @@ void cryptonight_core_gpu_hash(nvid_ctx* ctx, uint32_t nonce) } } -void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t startNonce, uint8_t version) +void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t startNonce) { if(miner_algo == cryptonight_monero) { - 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); + cryptonight_core_gpu_hash<CRYPTONIGHT_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight_monero>(ctx, startNonce); } else if(miner_algo == cryptonight_heavy) { - 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); - } + cryptonight_core_gpu_hash<CRYPTONIGHT_HEAVY_ITER, CRYPTONIGHT_HEAVY_MASK, CRYPTONIGHT_HEAVY_MEMORY/4, cryptonight_heavy>(ctx, startNonce); } else if(miner_algo == cryptonight) { @@ -502,5 +494,9 @@ void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t { cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_lite>(ctx, startNonce); } + else if(miner_algo == cryptonight_aeon) + { + cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_aeon>(ctx, startNonce); + } } diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu index c2a1f87..e2f0b2d 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu @@ -310,7 +310,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, xmrstak_algo miner_algo, uint8_t version) +extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce, xmrstak_algo miner_algo) { int threadsperblock = 128; uint32_t wsize = ctx->device_blocks * ctx->device_threads; @@ -318,7 +318,7 @@ extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce dim3 grid( ( wsize + threadsperblock - 1 ) / threadsperblock ); dim3 block( threadsperblock ); - if(miner_algo == cryptonight_heavy && version >= 3) + if(miner_algo == cryptonight_heavy) { 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 )); @@ -333,7 +333,7 @@ extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce } } -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) +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) { int threadsperblock = 128; uint32_t wsize = ctx->device_blocks * ctx->device_threads; @@ -344,7 +344,7 @@ 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 ) )); - if(miner_algo == cryptonight_heavy && version >= 3) + if(miner_algo == cryptonight_heavy) { CUDA_CHECK_MSG_KERNEL( ctx->device_id, diff --git a/xmrstak/cli/cli-miner.cpp b/xmrstak/cli/cli-miner.cpp index 2e84ec5..c425f04 100644 --- a/xmrstak/cli/cli-miner.cpp +++ b/xmrstak/cli/cli-miner.cpp @@ -79,6 +79,7 @@ void help() #endif #ifndef CONF_NO_OPENCL cout<<" --noAMD disable the AMD miner backend"<<endl; + cout<<" --noAMDCache disable the AMD(OpenCL) cache for precompiled binaries"<<endl; cout<<" --amd FILE AMD backend miner config file"<<endl; #endif #ifndef CONF_NO_CUDA @@ -449,6 +450,10 @@ int main(int argc, char *argv[]) { params::inst().useAMD = false; } + else if(opName.compare("--noAMDCache") == 0) + { + params::inst().AMDCache = false; + } else if(opName.compare("--noNVIDIA") == 0) { params::inst().useNVIDIA = false; diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp index 713beb4..f99698a 100644 --- a/xmrstak/jconf.cpp +++ b/xmrstak/jconf.cpp @@ -90,20 +90,23 @@ struct xmrstak_coin_algo { const char* coin_name; xmrstak_algo algo; + xmrstak_algo algo_root; + uint8_t fork_version; 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 }, - { "monero7", cryptonight_monero, "pool.usxmrpool.com:3333" }, - { "sumokoin", cryptonight_heavy, nullptr } + { "aeon7", cryptonight_aeon, cryptonight_lite, 7u, "mine.aeon-pool.com:5555" }, + { "cryptonight", cryptonight, cryptonight, 0u, nullptr }, + { "cryptonight_lite", cryptonight_lite, cryptonight_lite, 0u, nullptr }, + { "edollar", cryptonight, cryptonight, 0u, nullptr }, + { "electroneum", cryptonight, cryptonight, 0u, nullptr }, + { "graft", cryptonight, cryptonight, 0u, nullptr }, + { "intense", cryptonight, cryptonight, 0u, nullptr }, + { "karbo", cryptonight, cryptonight, 0u, nullptr }, + { "monero7", cryptonight_monero, cryptonight, 7u, "pool.usxmrpool.com:3333" }, + { "stellite", cryptonight_monero, cryptonight, 3u, nullptr }, + { "sumokoin", cryptonight_heavy, cryptonight, 3u, nullptr } }; constexpr size_t coin_alogo_size = (sizeof(coin_algos)/sizeof(coin_algos[0])); @@ -630,6 +633,8 @@ bool jconf::parse_config(const char* sFilename, const char* sFilenamePools) if(ctmp == coin_algos[i].coin_name) { mining_algo = coin_algos[i].algo; + mining_algo_root = coin_algos[i].algo_root; + mining_fork_version = coin_algos[i].fork_version; break; } } diff --git a/xmrstak/jconf.hpp b/xmrstak/jconf.hpp index 6874d37..76b93dd 100644 --- a/xmrstak/jconf.hpp +++ b/xmrstak/jconf.hpp @@ -48,7 +48,11 @@ public: bool TlsSecureAlgos(); - inline xmrstak_algo GetMiningAlgo() { return mining_algo; } + inline xmrstak_algo GetMiningAlgo() const { return mining_algo; } + + inline xmrstak_algo GetMiningAlgoRoot() const { return mining_algo_root; } + + inline uint8_t GetMiningForkVersion() const { return mining_fork_version; } std::string GetMiningCoin(); @@ -91,4 +95,6 @@ private: bool bHaveAes; xmrstak_algo mining_algo; + xmrstak_algo mining_algo_root; + uint8_t mining_fork_version; }; diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp index a620173..0e1dd9f 100644 --- a/xmrstak/misc/executor.cpp +++ b/xmrstak/misc/executor.cpp @@ -563,7 +563,7 @@ void executor::ex_main() else pools.emplace_front(0, "donate.xmr-stak.net:5500", "", "", "", 0.0, true, false, "", false); break; - + case cryptonight_aeon: case cryptonight_lite: if(dev_tls) pools.emplace_front(0, "donate.xmr-stak.net:7777", "", "", "", 0.0, true, true, "", true); diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp index 95bcc9c..d71aeb1 100644 --- a/xmrstak/net/jpsock.cpp +++ b/xmrstak/net/jpsock.cpp @@ -306,6 +306,12 @@ bool jpsock::process_line(char* line, size_t len) if(!mt->IsString()) return set_socket_error("PARSE error: Protocol error 1"); + if(strcmp(mt->GetString(), "mining.set_extranonce") == 0) + { + printer::inst()->print_msg(L0, "Detected buggy NiceHash pool code. Workaround engaged."); + return true; + } + if(strcmp(mt->GetString(), "job") != 0) return set_socket_error("PARSE error: Unsupported server method ", mt->GetString()); @@ -642,7 +648,13 @@ bool jpsock::cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bRes algo_name = "cryptonight-lite"; break; case cryptonight_monero: - algo_name = "cryptonight-monero"; + algo_name = "cryptonight-monerov7"; + break; + case cryptonight_aeon: + algo_name = "cryptonight-aeonv7"; + break; + case cryptonight_heavy: + algo_name = "cryptonight-heavy"; break; default: algo_name = "unknown"; diff --git a/xmrstak/params.hpp b/xmrstak/params.hpp index 6928df5..3d584b9 100644 --- a/xmrstak/params.hpp +++ b/xmrstak/params.hpp @@ -21,6 +21,7 @@ struct params std::string executablePrefix; std::string binaryName; bool useAMD; + bool AMDCache; bool useNVIDIA; bool useCPU; @@ -56,6 +57,7 @@ struct params binaryName("xmr-stak"), executablePrefix(""), useAMD(true), + AMDCache(true), useNVIDIA(true), useCPU(true), configFile("config.txt"), diff --git a/xmrstak/pools.tpl b/xmrstak/pools.tpl index d3a8a5a..f5afff6 100644 --- a/xmrstak/pools.tpl +++ b/xmrstak/pools.tpl @@ -20,7 +20,7 @@ POOLCONF], /* * Currency to mine. Supported values: * - * aeon + * aeon7 (use this for Aeon's new PoW) * cryptonight (try this if your coin is not listed) * cryptonight_lite * edollar diff --git a/xmrstak/version.cpp b/xmrstak/version.cpp index c7dbf9d..d3764e4 100644 --- a/xmrstak/version.cpp +++ b/xmrstak/version.cpp @@ -18,7 +18,7 @@ #endif #define XMR_STAK_NAME "xmr-stak" -#define XMR_STAK_VERSION "2.3.0" +#define XMR_STAK_VERSION "2.4.0" #if defined(_WIN32) #define OS_TYPE "win" |