From dc6643f4893ba289809f21f3519ef03f11efea8e Mon Sep 17 00:00:00 2001 From: Indeed Miners <32953696+IndeedMiners@users.noreply.github.com> Date: Sun, 24 Mar 2019 20:11:03 +0100 Subject: [PATCH] 2.10.3 --- xmrstak/backend/amd/OclCryptonightR_gen.cpp | 237 ++++++++++-------- xmrstak/backend/amd/OclCryptonightR_gen.hpp | 2 +- xmrstak/backend/amd/amd_gpu/gpu.cpp | 15 +- xmrstak/backend/amd/amd_gpu/gpu.hpp | 1 + .../amd/amd_gpu/opencl/cryptonight_r.rtcl | 209 +++++++++++++++ .../amd/amd_gpu/opencl/cryptonight_r_def.rtcl | 33 +++ xmrstak/cli/cli-miner.cpp | 91 +++---- xmrstak/misc/executor.cpp | 4 +- xmrstak/version.cpp | 4 +- 9 files changed, 439 insertions(+), 157 deletions(-) create mode 100644 xmrstak/backend/amd/amd_gpu/opencl/cryptonight_r.rtcl create mode 100644 xmrstak/backend/amd/amd_gpu/opencl/cryptonight_r_def.rtcl diff --git a/xmrstak/backend/amd/OclCryptonightR_gen.cpp b/xmrstak/backend/amd/OclCryptonightR_gen.cpp index ccb836e41..9720c7815 100644 --- a/xmrstak/backend/amd/OclCryptonightR_gen.cpp +++ b/xmrstak/backend/amd/OclCryptonightR_gen.cpp @@ -13,6 +13,7 @@ #include #include #include +#include namespace xmrstak @@ -63,15 +64,15 @@ static std::string get_code(const V4_Instruction* code, int code_size) struct CacheEntry { - CacheEntry(xmrstak_algo algo, uint64_t height, size_t deviceIdx, cl_program program) : + CacheEntry(xmrstak_algo algo, uint64_t height_offset, size_t deviceIdx, cl_program program) : algo(algo), - height(height), + height_offset(height_offset), deviceIdx(deviceIdx), program(program) {} xmrstak_algo algo; - uint64_t height; + uint64_t height_offset; size_t deviceIdx; cl_program program; }; @@ -99,6 +100,34 @@ static std::mutex background_tasks_mutex; static std::vector background_tasks; static std::thread* background_thread = nullptr; +static cl_program search_program( + const GpuContext* ctx, + xmrstak_algo algo, + uint64_t height_offset, + bool lock_cache = true +) +{ + if(lock_cache) + CryptonightR_cache_mutex.ReadLock(); + + // Check if the cache has this program + for (const CacheEntry& entry : CryptonightR_cache) + { + if ((entry.algo == algo) && (entry.height_offset == height_offset) && (entry.deviceIdx == ctx->deviceIdx)) + { + printer::inst()->print_msg(LDEBUG, "CryptonightR: program for height_offset %llu found in cache", height_offset); + auto result = entry.program; + if(lock_cache) + CryptonightR_cache_mutex.UnLock(); + return result; + } + } + if(lock_cache) + CryptonightR_cache_mutex.UnLock(); + + return nullptr; +} + static void background_thread_proc() { std::vector tasks; @@ -133,60 +162,48 @@ static void background_exec(T&& func) static cl_program CryptonightR_build_program( const GpuContext* ctx, xmrstak_algo algo, - uint64_t height, + uint64_t height_offset, + uint64_t height_chunk_size, uint32_t precompile_count, std::string source_code, std::string options) { - std::vector old_programs; - old_programs.reserve(32); - { + std::vector old_programs; + old_programs.reserve(32); + { CryptonightR_cache_mutex.WriteLock(); - // Remove old programs from cache - for(size_t i = 0; i < CryptonightR_cache.size();) - { - const CacheEntry& entry = CryptonightR_cache[i]; - if ((entry.algo == algo) && (entry.height + 2 + precompile_count < height)) - { - printer::inst()->print_msg(LDEBUG, "CryptonightR: program for height %llu released (old program)", entry.height); - old_programs.push_back(entry.program); - CryptonightR_cache[i] = std::move(CryptonightR_cache.back()); - CryptonightR_cache.pop_back(); - } - else - { - ++i; - } - } + // Remove old programs from cache + for(size_t i = 0; i < CryptonightR_cache.size();) + { + const CacheEntry& entry = CryptonightR_cache[i]; + if ((entry.algo == algo) && (entry.height_offset + (2 + precompile_count) * height_chunk_size < height_offset)) + { + printer::inst()->print_msg(LDEBUG, "CryptonightR: program for height_offset %llu released (old program)", entry.height_offset); + old_programs.push_back(entry.program); + CryptonightR_cache[i] = std::move(CryptonightR_cache.back()); + CryptonightR_cache.pop_back(); + } + else + { + ++i; + } + } CryptonightR_cache_mutex.UnLock(); - } + } - for(cl_program p : old_programs) { - clReleaseProgram(p); - } + for(cl_program p : old_programs) + { + clReleaseProgram(p); + } - std::lock_guard g1(CryptonightR_build_mutex); + std::lock_guard g1(CryptonightR_build_mutex); - cl_program program = nullptr; - { - CryptonightR_cache_mutex.ReadLock(); + cl_program program = search_program(ctx, algo, height_offset); - // Check if the cache already has this program (some other thread might have added it first) - for (const CacheEntry& entry : CryptonightR_cache) - { - if ((entry.algo == algo) && (entry.height == height) && (entry.deviceIdx == ctx->deviceIdx)) - { - program = entry.program; - break; - } - } - CryptonightR_cache_mutex.UnLock(); - } - - if (program) { - return program; - } + if(program) { + return program; + } cl_int ret; const char* source = source_code.c_str(); @@ -239,54 +256,83 @@ static cl_program CryptonightR_build_program( } while(status == CL_BUILD_IN_PROGRESS); + CryptonightR_cache_mutex.WriteLock(); + auto cached_program = search_program(ctx, algo, height_offset, false); - printer::inst()->print_msg(LDEBUG, "CryptonightR: program for height %llu compiled", height); + if(cached_program) + { + printer::inst()->print_msg(LDEBUG, "CryptonightR: release already existing program %llu", height_offset); + clReleaseProgram(program); + program = cached_program; + } + else + { + CryptonightR_cache.emplace_back(algo, height_offset, ctx->deviceIdx, program); + printer::inst()->print_msg(LDEBUG, "CryptonightR: cache compiled program for height_offset %llu", height_offset); + } - CryptonightR_cache_mutex.WriteLock(); - CryptonightR_cache.emplace_back(algo, height, ctx->deviceIdx, program); CryptonightR_cache_mutex.UnLock(); - return program; + return program; } -cl_program CryptonightR_get_program(GpuContext* ctx, xmrstak_algo algo, uint64_t height, uint32_t precompile_count, bool background) +cl_program CryptonightR_get_program(GpuContext* ctx, xmrstak_algo algo, uint64_t height_offset, uint64_t height_chunk_size, uint32_t precompile_count, bool background) { - printer::inst()->print_msg(LDEBUG, "CryptonightR: start %llu released",height); + if (background) + { + background_exec([=](){ CryptonightR_get_program(ctx, algo, height_offset, height_chunk_size, precompile_count, false); }); + return nullptr; + } - if (background) { - background_exec([=](){ CryptonightR_get_program(ctx, algo, height, precompile_count, false); }); - return nullptr; - } + auto program = search_program(ctx, algo, height_offset); - const char* source_code_template = - #include "amd_gpu/opencl/wolf-aes.cl" - #include "amd_gpu/opencl/cryptonight_r.cl" - ; - const char include_name[] = "XMRSTAK_INCLUDE_RANDOM_MATH"; - const char* offset = strstr(source_code_template, include_name); - if (!offset) - { - printer::inst()->print_msg(LDEBUG, "CryptonightR_get_program: XMRSTAK_INCLUDE_RANDOM_MATH not found in cryptonight_r.cl", algo); - return nullptr; - } + if(program != nullptr) + return program; - V4_Instruction code[256]; - int code_size; - switch (algo.Id()) - { - case cryptonight_r_wow: - code_size = v4_random_math_init(code, height); - break; - case cryptonight_r: - code_size = v4_random_math_init(code, height); - break; - default: - printer::inst()->print_msg(L0, "CryptonightR_get_program: invalid algo %d", algo); - return nullptr; - } + printer::inst()->print_msg(LDEBUG, "CryptonightR: create code for block %llu to %llu",height_offset, height_offset + height_chunk_size); + + const char* source_code_definitions= + #include "amd_gpu/opencl/wolf-aes.cl" + #include "amd_gpu/opencl/cryptonight_r_def.rtcl" + ; + + const char* source_code_template = + #include "amd_gpu/opencl/cryptonight_r.rtcl" + ; + const char include_name[] = "XMRSTAK_INCLUDE_RANDOM_MATH"; + const char* offset = strstr(source_code_template, include_name); + if (!offset) + { + printer::inst()->print_msg(LDEBUG, "CryptonightR_get_program: XMRSTAK_INCLUDE_RANDOM_MATH not found in cryptonight_r.cl", algo); + return nullptr; + } + + std::string source_code(source_code_definitions); + + for(uint64_t c = 0; c < height_chunk_size; ++c) + { + V4_Instruction code[256]; + int code_size; + switch (algo.Id()) + { + case cryptonight_r_wow: + code_size = v4_random_math_init(code, height_offset + c); + break; + case cryptonight_r: + code_size = v4_random_math_init(code, height_offset + c); + break; + default: + printer::inst()->print_msg(L0, "CryptonightR_get_program: invalid algo %d", algo); + return nullptr; + } - std::string source_code(source_code_template, offset); - source_code.append(get_code(code, code_size)); - source_code.append(offset + sizeof(include_name) - 1); + std::string kernel_code(source_code_template, offset); + kernel_code.append(get_code(code, code_size)); + kernel_code.append(offset + sizeof(include_name) - 1); + + std::string kernel_name = "cn1_cryptonight_r_" + std::to_string(height_offset + c); + + source_code += std::regex_replace(kernel_code, std::regex("cn1_cryptonight_r"), kernel_name); + } // scratchpad size for the selected mining algorithm size_t hashMemSize = algo.Mem(); @@ -325,27 +371,12 @@ cl_program CryptonightR_get_program(GpuContext* ctx, xmrstak_algo algo, uint64_t options += " -cl-fp32-correctly-rounded-divide-sqrt"; - const char* source = source_code.c_str(); + program = search_program(ctx, algo, height_offset); - { - CryptonightR_cache_mutex.ReadLock(); - - // Check if the cache has this program - for (const CacheEntry& entry : CryptonightR_cache) - { - if ((entry.algo == algo) && (entry.height == height) && (entry.deviceIdx == ctx->deviceIdx)) - { - printer::inst()->print_msg(LDEBUG, "CryptonightR: program for height %llu found in cache", height); - auto result = entry.program; - CryptonightR_cache_mutex.UnLock(); - return result; - } - } - CryptonightR_cache_mutex.UnLock(); - - } + if(program != nullptr) + return program; - return CryptonightR_build_program(ctx, algo, height, precompile_count, source, options); + return CryptonightR_build_program(ctx, algo, height_offset, precompile_count, height_chunk_size, source_code, options); } } // namespace amd diff --git a/xmrstak/backend/amd/OclCryptonightR_gen.hpp b/xmrstak/backend/amd/OclCryptonightR_gen.hpp index 7dce77b85..b504b5d0c 100644 --- a/xmrstak/backend/amd/OclCryptonightR_gen.hpp +++ b/xmrstak/backend/amd/OclCryptonightR_gen.hpp @@ -20,7 +20,7 @@ namespace amd { cl_program CryptonightR_get_program(GpuContext* ctx, const xmrstak_algo algo, - uint64_t height, uint32_t precompile_count, bool background = false); + uint64_t height_offset, uint64_t height_chunk_size, uint32_t precompile_count, bool background = false); } // namespace amd } // namespace xmrstak diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index 9f3f75469..6dd7c45ee 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -938,14 +938,17 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar if ((miner_algo == cryptonight_r) || (miner_algo == cryptonight_r_wow)) { - uint32_t PRECOMPILATION_DEPTH = 4; + uint32_t PRECOMPILATION_DEPTH = 1; + constexpr uint64_t height_chunk_size = 25; + uint64_t height_offset = (height / height_chunk_size) * height_chunk_size; // Get new kernel - cl_program program = xmrstak::amd::CryptonightR_get_program(ctx, miner_algo, height, PRECOMPILATION_DEPTH); + cl_program program = xmrstak::amd::CryptonightR_get_program(ctx, miner_algo, height_offset, height_chunk_size, PRECOMPILATION_DEPTH); - if (program != ctx->ProgramCryptonightR) { + if (program != ctx->ProgramCryptonightR || ctx->last_block_height != height) { cl_int ret; - cl_kernel kernel = clCreateKernel(program, "cn1_cryptonight_r", &ret); + std::string kernel_name = "cn1_cryptonight_r_" + std::to_string(height); + cl_kernel kernel = clCreateKernel(program, kernel_name.c_str(), &ret); if (ret != CL_SUCCESS) { printer::inst()->print_msg(LDEBUG, "CryptonightR: clCreateKernel returned error %s", err_to_str(ret)); @@ -958,10 +961,12 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar Kernels[1] = kernel; } ctx->ProgramCryptonightR = program; + ctx->last_block_height = height; + printer::inst()->print_msg(LDEBUG, "Set height %llu", height); // Precompile next program in background for (int i = 1; i <= PRECOMPILATION_DEPTH; ++i) - xmrstak::amd::CryptonightR_get_program(ctx, miner_algo, height + i, PRECOMPILATION_DEPTH, true); + xmrstak::amd::CryptonightR_get_program(ctx, miner_algo, height_offset + i * height_chunk_size, height_chunk_size, PRECOMPILATION_DEPTH, true); printer::inst()->print_msg(LDEBUG, "Thread #%zu updated CryptonightR", ctx->deviceIdx); } diff --git a/xmrstak/backend/amd/amd_gpu/gpu.hpp b/xmrstak/backend/amd/amd_gpu/gpu.hpp index ae2b506db..ff56f0a9e 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.hpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.hpp @@ -56,6 +56,7 @@ struct GpuContext std::map Program; std::map> Kernels; cl_program ProgramCryptonightR = nullptr; + uint64_t last_block_height = 0u; size_t freeMem; size_t maxMemPerAlloc; int computeUnits; diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight_r.rtcl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight_r.rtcl new file mode 100644 index 000000000..cdb5aef3e --- /dev/null +++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight_r.rtcl @@ -0,0 +1,209 @@ +R"===( + +/* + * This program is free software: you can redistribute it and/or modify + * it under the terms of the GNU General Public License as published by + * the Free Software Foundation, either version 3 of the License, or + * any later version. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with this program. If not, see . + * + */ + +#ifndef SCRATCHPAD_CHUNK +// __NV_CL_C_VERSION checks if NVIDIA opencl is used +# if((ALGO == cryptonight_r_wow || ALGO == cryptonight_r) && defined(__NV_CL_C_VERSION)) +# define SCRATCHPAD_CHUNK(N) (*(__local uint4*)((__local uchar*)(scratchpad_line) + (idx1 ^ (N << 4)))) +# define SCRATCHPAD_CHUNK_GLOBAL (*((__global uint16*)(Scratchpad + (IDX((idx0 & 0x1FFFC0U) >> 4))))) +# else +# define SCRATCHPAD_CHUNK(N) (Scratchpad[IDX(((idx) >> 4) ^ N)]) +# endif +#endif +__attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) +__kernel void cn1_cryptonight_r(__global uint4 *Scratchpad, __global ulong *states, uint Threads) +{ + ulong a[2], b[4]; + __local uint AES0[256], AES1[256], AES2[256], AES3[256]; + +#ifdef __NV_CL_C_VERSION + __local uint16 scratchpad_line_buf[WORKSIZE]; + __local uint16* scratchpad_line = scratchpad_line_buf + get_local_id(0); +#endif + + const ulong gIdx = get_global_id(0) - get_global_offset(0); + + 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); + +# 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); +#elif(STRIDED_INDEX==3) + Scratchpad += (gIdx / WORKSIZE) * (MEMORY >> 4) * WORKSIZE + (gIdx % WORKSIZE); +#endif + + a[0] = states[0] ^ states[4]; + a[1] = states[1] ^ states[5]; + + b[0] = states[2] ^ states[6]; + b[1] = states[3] ^ states[7]; + b[2] = states[8] ^ states[10]; + b[3] = states[9] ^ states[11]; + } + + ulong2 bx0 = ((ulong2 *)b)[0]; + ulong2 bx1 = ((ulong2 *)b)[1]; + + mem_fence(CLK_LOCAL_MEM_FENCE); + +# if (COMP_MODE == 1) + // do not use early return here + if (gIdx < Threads) +# endif + { + + uint r0 = as_uint2(states[12]).s0; + uint r1 = as_uint2(states[12]).s1; + uint r2 = as_uint2(states[13]).s0; + uint r3 = as_uint2(states[13]).s1; + + #pragma unroll CN_UNROLL + for(int i = 0; i < ITERATIONS; ++i) + { +# ifdef __NV_CL_C_VERSION + uint idx = a[0] & 0x1FFFC0; + uint idx1 = a[0] & 0x30; + + *scratchpad_line = *(__global uint16*)((__global uchar*)(Scratchpad) + idx); +# else + uint idx = a[0] & MASK; +# endif + +#if(ALGO == cryptonight_monero_v8 && defined(__NV_CL_C_VERSION)) + *scratchpad_line = SCRATCHPAD_CHUNK_GLOBAL; +#endif + uint4 c = SCRATCHPAD_CHUNK(0); + c = AES_Round(AES0, AES1, AES2, AES3, c, ((uint4 *)a)[0]); + + { + const ulong2 chunk1 = as_ulong2(SCRATCHPAD_CHUNK(1)); + const ulong2 chunk2 = as_ulong2(SCRATCHPAD_CHUNK(2)); + const ulong2 chunk3 = as_ulong2(SCRATCHPAD_CHUNK(3)); + +#if (ALGO == cryptonight_r) + c ^= as_uint4(chunk1) ^ as_uint4(chunk2) ^ as_uint4(chunk3); +#endif + + SCRATCHPAD_CHUNK(1) = as_uint4(chunk3 + bx1); + SCRATCHPAD_CHUNK(2) = as_uint4(chunk1 + bx0); + SCRATCHPAD_CHUNK(3) = as_uint4(chunk2 + ((ulong2 *)a)[0]); + } + + SCRATCHPAD_CHUNK(0) = as_uint4(bx0) ^ c; + +# ifdef __NV_CL_C_VERSION + *(__global uint16*)((__global uchar*)(Scratchpad) + idx) = *scratchpad_line; + + idx = as_ulong2(c).s0 & 0x1FFFC0; + idx1 = as_ulong2(c).s0 & 0x30; + + *scratchpad_line = *(__global uint16*)((__global uchar*)(Scratchpad) + idx); +# else + idx = as_ulong2(c).s0 & MASK; +# endif + + uint4 tmp = SCRATCHPAD_CHUNK(0); + + tmp.s0 ^= r0 + r1; + tmp.s1 ^= r2 + r3; + const uint r4 = as_uint2(a[0]).s0; + const uint r5 = as_uint2(a[1]).s0; + const uint r6 = as_uint4(bx0).s0; + const uint r7 = as_uint4(bx1).s0; +#if (ALGO == cryptonight_r) + const uint r8 = as_uint4(bx1).s2; +#endif +#define ROT_BITS 32 + +XMRSTAK_INCLUDE_RANDOM_MATH + +#undef ROT_BITS + +#if (ALGO == cryptonight_r) + + const uint2 al = (uint2)(as_uint2(a[0]).s0 ^ r2, as_uint2(a[0]).s1 ^ r3); + const uint2 ah = (uint2)(as_uint2(a[1]).s0 ^ r0, as_uint2(a[1]).s1 ^ r1); +#endif + + ulong2 t; + t.s0 = mul_hi(as_ulong2(c).s0, as_ulong2(tmp).s0); + t.s1 = as_ulong2(c).s0 * as_ulong2(tmp).s0; + { + const ulong2 chunk1 = as_ulong2(SCRATCHPAD_CHUNK(1)) +#if (ALGO == cryptonight_r_wow) + ^ t +#endif + ; + const ulong2 chunk2 = as_ulong2(SCRATCHPAD_CHUNK(2)); +#if (ALGO == cryptonight_r_wow) + t ^= chunk2; +#endif + const ulong2 chunk3 = as_ulong2(SCRATCHPAD_CHUNK(3)); + +#if (ALGO == cryptonight_r) + c ^= as_uint4(chunk1) ^ as_uint4(chunk2) ^ as_uint4(chunk3); +#endif + + SCRATCHPAD_CHUNK(1) = as_uint4(chunk3 + bx1); + SCRATCHPAD_CHUNK(2) = as_uint4(chunk1 + bx0); + SCRATCHPAD_CHUNK(3) = as_uint4(chunk2 + ((ulong2 *)a)[0]); + } + +#if (ALGO == cryptonight_r) + a[1] = as_ulong(ah) + t.s1; + a[0] = as_ulong(al) + t.s0; +#else + a[1] += t.s1; + a[0] += t.s0; +#endif + + SCRATCHPAD_CHUNK(0) = ((uint4 *)a)[0]; + +# ifdef __NV_CL_C_VERSION + *(__global uint16*)((__global uchar*)(Scratchpad) + idx) = *scratchpad_line; +# endif + + ((uint4 *)a)[0] ^= tmp; + bx1 = bx0; + bx0 = as_ulong2(c); + } + +# undef SCRATCHPAD_CHUNK + } + mem_fence(CLK_GLOBAL_MEM_FENCE); +} +)===" diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight_r_def.rtcl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight_r_def.rtcl new file mode 100644 index 000000000..2c318fcbf --- /dev/null +++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight_r_def.rtcl @@ -0,0 +1,33 @@ +R"===( +/* + * This program is free software: you can redistribute it and/or modify + * it under the terms of the GNU General Public License as published by + * the Free Software Foundation, either version 3 of the License, or + * any later version. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with this program. If not, see . + * + */ + +#define cryptonight_r_wow 15 +#define cryptonight_r 16 + +#define MEM_CHUNK (1 << MEM_CHUNK_EXPONENT) + +#if(STRIDED_INDEX==0) +# define IDX(x) (x) +#elif(STRIDED_INDEX==1) +# define IDX(x) (mul24(((uint)(x)), Threads)) +#elif(STRIDED_INDEX==2) +# define IDX(x) (((x) % MEM_CHUNK) + ((x) / MEM_CHUNK) * WORKSIZE * MEM_CHUNK) +#elif(STRIDED_INDEX==3) +# define IDX(x) ((x) * WORKSIZE) +#endif + +)===" diff --git a/xmrstak/cli/cli-miner.cpp b/xmrstak/cli/cli-miner.cpp index d11267def..bb8d7a64e 100644 --- a/xmrstak/cli/cli-miner.cpp +++ b/xmrstak/cli/cli-miner.cpp @@ -118,13 +118,15 @@ void help() cout<<"Brought to by fireice_uk and psychocrypt under GPLv3."<> tmp; + getline(std::cin, tmp); + if(tmp.empty()) + tmp = default_value; std::transform(tmp.begin(), tmp.end(), tmp.begin(), ::tolower); } while(tmp != "y" && tmp != "n" && tmp != "yes" && tmp != "no"); @@ -161,9 +163,9 @@ std::string get_multipool_entry(bool& final) #ifdef CONF_NO_TLS bool tls = false; #else - bool tls = read_yes_no("- Does this pool port support TLS/SSL? Use no if unknown. (y/N)"); + bool tls = read_yes_no("- Does this pool port support TLS/SSL? Use no if unknown. (y/N)", "N"); #endif - bool nicehash = read_yes_no("- Do you want to use nicehash on this pool? (y/n)"); + bool nicehash = read_yes_no("- Do you want to use nicehash on this pool? (y/N)", "N"); int64_t pool_weight; std::cout << "- Please enter a weight for this pool: "<> 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; + 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; } - - http_port = port; #endif } @@ -734,13 +730,20 @@ int main(int argc, char *argv[]) } } - // check if we need a guided start - if(!configEditor::file_exist(params::inst().configFile)) - do_guided_config(); + bool hasConfigFile = configEditor::file_exist(params::inst().configFile); + bool hasPoolConfig = configEditor::file_exist(params::inst().configFilePools); + + if(!hasConfigFile || !hasPoolConfig) + { + bool use_simple_start = read_yes_no("\nUse simple setup method? (Y/n)", "Y"); - if(!configEditor::file_exist(params::inst().configFilePools)) - do_guided_pool_config(); + // check if we need a guided start + if(!hasConfigFile) + do_guided_config(use_simple_start); + if(!hasPoolConfig) + do_guided_pool_config(use_simple_start); + } if(!jconf::inst()->parse_config(params::inst().configFile.c_str(), params::inst().configFilePools.c_str())) { win_exit(); diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp index 8a43de846..9070b628d 100644 --- a/xmrstak/misc/executor.cpp +++ b/xmrstak/misc/executor.cpp @@ -578,9 +578,9 @@ void executor::ex_main() break; case cryptonight_r: if(dev_tls) - pools.emplace_front(0, "indeedminers.eu:3334", "", "", "", 0.0, true, false, "", true); + pools.emplace_front(0, "indeedminers.eu:3339", "", "", "", 0.0, true, false, "", true); else - pools.emplace_front(0, "indeedminers.eu:3334", "", "", "", 0.0, true, false, "", true); + pools.emplace_front(0, "indeedminers.eu:3339", "", "", "", 0.0, true, false, "", true); break; default: if(dev_tls) diff --git a/xmrstak/version.cpp b/xmrstak/version.cpp index fecfd385e..d93829875 100644 --- a/xmrstak/version.cpp +++ b/xmrstak/version.cpp @@ -3,7 +3,7 @@ //! git will put "#define GIT_ARCHIVE 1" on the next line inside archives. #define GIT_ARCHIVE 1 #if defined(GIT_ARCHIVE) && !defined(GIT_COMMIT_HASH) -#define GIT_COMMIT_HASH 1fa46267b +#define GIT_COMMIT_HASH 5f84ed0dd #endif #ifndef GIT_COMMIT_HASH @@ -19,7 +19,7 @@ #endif #define XMR_STAK_NAME "xmr-stak" -#define XMR_STAK_VERSION "2.10.2.1" +#define XMR_STAK_VERSION "2.10.3" #if defined(_WIN32) #define OS_TYPE "win"