diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index 43fa12c1a..03100d0a3 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -309,8 +309,8 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_ } size_t scratchPadSize = std::max( - cn_select_memory(::jconf::inst()->GetMiningAlgo()), - cn_select_memory(::jconf::inst()->GetMiningAlgoRoot()) + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()), + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot()) ); size_t g_thd = ctx->rawIntensity; @@ -376,8 +376,8 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_ } xmrstak_algo miner_algo[2] = { - ::jconf::inst()->GetMiningAlgo(), - ::jconf::inst()->GetMiningAlgoRoot() + ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo(), + ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot() }; int num_algos = miner_algo[0] == miner_algo[1] ? 1 : 2; @@ -675,11 +675,18 @@ std::vector getAMDDevices(int index) } std::string devVendor(devVendorVec.data()); - if( devVendor.find("Advanced Micro Devices") != std::string::npos || devVendor.find("AMD") != std::string::npos) + + bool isAMDDevice = devVendor.find("Advanced Micro Devices") != std::string::npos || devVendor.find("AMD") != std::string::npos; + bool isNVIDIADevice = devVendor.find("NVIDIA Corporation") != std::string::npos || devVendor.find("NVIDIA") != std::string::npos; + + std::string selectedOpenCLVendor = xmrstak::params::inst().openCLVendor; + if((isAMDDevice && selectedOpenCLVendor == "AMD") || (isNVIDIADevice && selectedOpenCLVendor == "NVIDIA")) { GpuContext ctx; std::vector devNameVec(1024); size_t maxMem; + if( devVendor.find("NVIDIA Corporation") != std::string::npos) + ctx.isNVIDIA = true; if((clStatus = clGetDeviceInfo(device_list[k], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(int), &(ctx.computeUnits), NULL)) != CL_SUCCESS) { @@ -699,6 +706,10 @@ std::vector getAMDDevices(int index) continue; } + // the allocation for NVIDIA OpenCL is not limited to 1/4 of the GPU memory per allocation + if(ctx.isNVIDIA) + maxMem = ctx.freeMem; + if((clStatus = clGetDeviceInfo(device_list[k], CL_DEVICE_NAME, devNameVec.size(), devNameVec.data(), NULL)) != CL_SUCCESS) { printer::inst()->print_msg(L1,"WARNING: %s when calling clGetDeviceInfo to get CL_DEVICE_NAME for device %u.", err_to_str(clStatus), k); @@ -747,13 +758,15 @@ int getAMDPlatformIdx() clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, infoSize, platformNameVec.data(), NULL); std::string platformName(platformNameVec.data()); - if( platformName.find("Advanced Micro Devices") != std::string::npos || + + bool isAMDOpenCL = platformName.find("Advanced Micro Devices") != std::string::npos || platformName.find("Apple") != std::string::npos || - platformName.find("Mesa") != std::string::npos - ) + platformName.find("Mesa") != std::string::npos; + bool isNVIDIADevice = platformName.find("NVIDIA Corporation") != std::string::npos || platformName.find("NVIDIA") != std::string::npos; + std::string selectedOpenCLVendor = xmrstak::params::inst().openCLVendor; + if((isAMDOpenCL && selectedOpenCLVendor == "AMD") || (isNVIDIADevice && selectedOpenCLVendor == "NVIDIA")) { - - printer::inst()->print_msg(L0,"Found AMD platform index id = %i, name = %s",i , platformName.c_str()); + printer::inst()->print_msg(L0,"Found %s platform index id = %i, name = %s", selectedOpenCLVendor.c_str(), i , platformName.c_str()); if(platformName.find("Mesa") != std::string::npos) mesaPlatform = i; else @@ -819,7 +832,7 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx) std::vector platformNameVec(infoSize); clGetPlatformInfo(PlatformIDList[platform_idx], CL_PLATFORM_VENDOR, infoSize, platformNameVec.data(), NULL); std::string platformName(platformNameVec.data()); - if( platformName.find("Advanced Micro Devices") == std::string::npos) + if(xmrstak::params::inst().openCLVendor == "AMD" && platformName.find("Advanced Micro Devices") == std::string::npos) { printer::inst()->print_msg(L1,"WARNING: using non AMD device: %s", platformName.c_str()); } @@ -907,7 +920,8 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx) { size_t reduced_intensity = (ctx[i].rawIntensity / ctx[i].workSize) * ctx[i].workSize; ctx[i].rawIntensity = reduced_intensity; - printer::inst()->print_msg(L0, "WARNING AMD: gpu %d intensity is not a multiple of 'worksize', auto reduce intensity to %d", ctx[i].deviceIdx, int(reduced_intensity)); + const std::string backendName = xmrstak::params::inst().openCLVendor; + printer::inst()->print_msg(L0, "WARNING %s: gpu %d intensity is not a multiple of 'worksize', auto reduce intensity to %d", backendName.c_str(), ctx[i].deviceIdx, int(reduced_intensity)); } if((ret = InitOpenCLGpu(opencl_ctx, &ctx[i], source_code.c_str())) != ERR_SUCCESS) @@ -922,7 +936,7 @@ 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, xmrstak_algo miner_algo) { // switch to the kernel storage - int kernel_storage = miner_algo == ::jconf::inst()->GetMiningAlgo() ? 0 : 1; + int kernel_storage = miner_algo == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() ? 0 : 1; cl_int ret; @@ -990,7 +1004,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar return(ERR_OCL_API); } - if(miner_algo == cryptonight_ipbc || miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon ) + if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon ) { // Input if ((ret = clSetKernelArg(ctx->Kernels[kernel_storage][1], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS) @@ -1087,7 +1101,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo) { // switch to the kernel storage - int kernel_storage = miner_algo == ::jconf::inst()->GetMiningAlgo() ? 0 : 1; + int kernel_storage = miner_algo == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() ? 0 : 1; cl_int ret; cl_uint zero = 0; diff --git a/xmrstak/backend/amd/amd_gpu/gpu.hpp b/xmrstak/backend/amd/amd_gpu/gpu.hpp index 0db6c9076..5ab80b82a 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.hpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.hpp @@ -27,6 +27,7 @@ struct GpuContext size_t workSize; int stridedIndex; int memChunk; + bool isNVIDIA = false; int compMode; /*Output vars*/ diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl index 977efb452..d2ae1a7ee 100644 --- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl +++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl @@ -32,8 +32,7 @@ R"===( * The implemented function is modified because the last is in our case always a scalar. * We can ignore the bitwise AND operation. */ - -inline uint2 __attribute__((overloadable)) amd_bitalign( const uint2 src0, const uint2 src1, const uint src2) +inline uint2 amd_bitalign( const uint2 src0, const uint2 src1, const uint src2) { uint2 result; result.s0 = (uint) (((((long)src0.s0) << 32) | (long)src1.s0) >> (src2)); @@ -60,7 +59,7 @@ inline uint2 __attribute__((overloadable)) amd_bitalign( const uint2 src0, const * dst.s0 = src0.s0 >> offset; * similar operation applied to other components of the vectors */ -inline uint __attribute__((overloadable)) amd_bfe(const uint src0, const uint offset, const uint width) +inline int amd_bfe(const uint src0, const uint offset, const uint width) { /* casts are removed because we can implement everything as uint * int offset = src1; @@ -554,8 +553,8 @@ __kernel void JOIN(cn0,ALGO)(__global ulong *input, __global uint4 *Scratchpad, __attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states, ulong Threads -// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc -#if(ALGO == 3 || ALGO == 5 || ALGO == 6) +// cryptonight_monero || cryptonight_aeon +#if(ALGO == 3 || ALGO == 5) , __global ulong *input #endif ) @@ -575,7 +574,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states } barrier(CLK_LOCAL_MEM_FENCE); -#if(ALGO == 3 || ALGO == 5 || ALGO == 6) +#if(ALGO == 3 || ALGO == 5) uint2 tweak1_2; #endif uint4 b_x; @@ -599,7 +598,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states b[1] = states[3] ^ states[7]; b_x = ((uint4 *)b)[0]; -#if(ALGO == 3 || ALGO == 5 || ALGO == 6) +#if(ALGO == 3 || ALGO == 5) tweak1_2 = as_uint2(input[4]); tweak1_2.s0 >>= 24; tweak1_2.s0 |= tweak1_2.s1 << 8; @@ -626,7 +625,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states ((uint4 *)c)[0] = AES_Round(AES0, AES1, AES2, AES3, ((uint4 *)c)[0], ((uint4 *)a)[0]); b_x ^= ((uint4 *)c)[0]; -#if(ALGO == 3 || ALGO == 5 || ALGO == 6) +#if(ALGO == 3 || ALGO == 5) uint table = 0x75310U; uint index = ((b_x.s2 >> 26) & 12) | ((b_x.s2 >> 23) & 2); b_x.s2 ^= ((table >> index) & 0x30U) << 24; @@ -640,18 +639,13 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states a[0] += mul_hi(c[0], as_ulong2(tmp).s0); -#if(ALGO == 3 || ALGO == 5 || ALGO == 6) +#if(ALGO == 3 || ALGO == 5) ((uint2 *)&(a[1]))[0] ^= tweak1_2; Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0]; ((uint2 *)&(a[1]))[0] ^= tweak1_2; #else Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0]; #endif -#if(ALGO == 6) - long prev = *((__global long*)(Scratchpad + (IDX((c[0] & MASK) >> 4)))); - long cur = *((__global long*)(Scratchpad + (IDX((c[0] & MASK) >> 4))) + 1); - *((__global long*)(Scratchpad + (IDX((c[0] & MASK) >> 4))) + 1) = prev ^ cur; -#endif ((uint4 *)a)[0] ^= tmp; idx0 = a[0]; diff --git a/xmrstak/backend/amd/autoAdjust.hpp b/xmrstak/backend/amd/autoAdjust.hpp index 6df0eeaa0..685890bc2 100644 --- a/xmrstak/backend/amd/autoAdjust.hpp +++ b/xmrstak/backend/amd/autoAdjust.hpp @@ -84,13 +84,14 @@ class autoAdjust constexpr size_t byteToMiB = 1024u * 1024u; size_t hashMemSize = std::max( - cn_select_memory(::jconf::inst()->GetMiningAlgo()), - cn_select_memory(::jconf::inst()->GetMiningAlgoRoot()) + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()), + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot()) ); std::string conf; for(auto& ctx : devVec) { + size_t minFreeMem = 128u * byteToMiB; /* 1000 is a magic selected limit, the reason is that more than 2GiB memory * sowing down the memory performance because of TLB cache misses */ @@ -112,12 +113,26 @@ class autoAdjust */ maxThreads = 2024u; } + + // NVIDIA optimizations + if( + ctx.isNVIDIA && ( + ctx.name.find("P100") != std::string::npos || + ctx.name.find("V100") != std::string::npos + ) + ) + { + // do not limit the number of threads + maxThreads = 40000u; + minFreeMem = 512u * byteToMiB; + } + // increase all intensity limits by two for aeon - if(::jconf::inst()->GetMiningAlgo() == cryptonight_lite) + if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_lite) maxThreads *= 2u; // keep 128MiB memory free (value is randomly chosen) - size_t availableMem = ctx.freeMem - (128u * byteToMiB); + size_t availableMem = ctx.freeMem - minFreeMem; // 224byte extra memory is used per thread for meta data size_t perThread = hashMemSize + 224u; size_t maxIntensity = availableMem / perThread; @@ -138,7 +153,7 @@ class autoAdjust // set 8 threads per block (this is a good value for the most gpus) conf += std::string(" { \"index\" : ") + std::to_string(ctx.deviceIdx) + ",\n" + " \"intensity\" : " + std::to_string(intensity) + ", \"worksize\" : " + std::to_string(8) + ",\n" + - " \"affine_to_cpu\" : false, \"strided_index\" : 1, \"mem_chunk\" : 2,\n" + " \"affine_to_cpu\" : false, \"strided_index\" : " + (ctx.isNVIDIA ? "0" : "1") + ", \"mem_chunk\" : 2,\n" " \"comp_mode\" : true\n" + " },\n"; } @@ -151,7 +166,9 @@ class autoAdjust configTpl.replace("PLATFORMINDEX",std::to_string(platformIndex)); configTpl.replace("GPUCONFIG",conf); configTpl.write(params::inst().configFileAMD); - printer::inst()->print_msg(L0, "AMD: GPU configuration stored in file '%s'", params::inst().configFileAMD.c_str()); + + const std::string backendName = xmrstak::params::inst().openCLVendor; + printer::inst()->print_msg(L0, "%s: GPU (OpenCL) configuration stored in file '%s'", backendName.c_str(), params::inst().configFileAMD.c_str()); } std::vector devVec; diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp index 76c6676c6..4353e3d05 100644 --- a/xmrstak/backend/amd/minethd.cpp +++ b/xmrstak/backend/amd/minethd.cpp @@ -137,6 +137,8 @@ std::vector* minethd::thread_starter(uint32_t threadOffset, miner_wor for (i = 0; i < n; i++) { jconf::inst()->GetThreadConfig(i, cfg); + + const std::string backendName = xmrstak::params::inst().openCLVendor; if(cfg.cpu_aff >= 0) { @@ -144,10 +146,10 @@ std::vector* minethd::thread_starter(uint32_t threadOffset, miner_wor printer::inst()->print_msg(L1, "WARNING on macOS thread affinity is only advisory."); #endif - printer::inst()->print_msg(L1, "Starting AMD GPU thread %d, affinity: %d.", i, (int)cfg.cpu_aff); + printer::inst()->print_msg(L1, "Starting %s GPU (OpenCL) thread %d, affinity: %d.", backendName.c_str(), i, (int)cfg.cpu_aff); } else - printer::inst()->print_msg(L1, "Starting AMD GPU thread %d, no affinity.", i); + printer::inst()->print_msg(L1, "Starting %s GPU (OpenCL) thread %d, no affinity.", backendName.c_str(), i); minethd* thd = new minethd(pWork, i + threadOffset, &vGpuData[i], cfg); pvThreads->push_back(thd); @@ -193,12 +195,13 @@ void minethd::work_main() cpu_ctx = cpu::minethd::minethd_alloc_ctx(); // start with root algorithm and switch later if fork version is reached - auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot(); + auto miner_algo = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot(); cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); globalStates::inst().iConsumeCnt++; uint8_t version = 0; + size_t lastPoolId = 0; while (bQuit == 0) { @@ -217,14 +220,20 @@ void minethd::work_main() } uint8_t new_version = oWork.getVersion(); - if (::jconf::inst()->GetMiningAlgo() == cryptonight_ipbc) new_version = oWork.bWorkBlob[1]; - if(new_version != version) + if(new_version != version || oWork.iPoolId != lastPoolId) { - if(new_version >= ::jconf::inst()->GetMiningForkVersion()) + coinDescription coinDesc = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(oWork.iPoolId); + if(new_version >= coinDesc.GetMiningForkVersion()) + { + miner_algo = coinDesc.GetMiningAlgo(); + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); + } + else { - miner_algo = ::jconf::inst()->GetMiningAlgo(); + miner_algo = coinDesc.GetMiningAlgoRoot(); hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); } + lastPoolId = oWork.iPoolId; version = new_version; } diff --git a/xmrstak/backend/backendConnector.cpp b/xmrstak/backend/backendConnector.cpp index acedbd6a5..6f80a0f73 100644 --- a/xmrstak/backend/backendConnector.cpp +++ b/xmrstak/backend/backendConnector.cpp @@ -77,11 +77,12 @@ std::vector* BackendConnector::thread_starter(miner_work& pWork) #ifndef CONF_NO_OPENCL if(params::inst().useAMD) { - plugin amdplugin("AMD", "xmrstak_opencl_backend"); + const std::string backendName = xmrstak::params::inst().openCLVendor; + plugin amdplugin(backendName, "xmrstak_opencl_backend"); std::vector* amdThreads = amdplugin.startBackend(static_cast(pvThreads->size()), pWork, environment::inst()); pvThreads->insert(std::end(*pvThreads), std::begin(*amdThreads), std::end(*amdThreads)); if(amdThreads->size() == 0) - printer::inst()->print_msg(L0, "WARNING: backend AMD disabled."); + printer::inst()->print_msg(L0, "WARNING: backend %s (OpenCL) disabled.", backendName.c_str()); } #endif diff --git a/xmrstak/backend/cpu/autoAdjust.hpp b/xmrstak/backend/cpu/autoAdjust.hpp index ed96d8b7b..518721a2f 100644 --- a/xmrstak/backend/cpu/autoAdjust.hpp +++ b/xmrstak/backend/cpu/autoAdjust.hpp @@ -37,8 +37,8 @@ class autoAdjust { const size_t hashMemSizeKB = std::max( - cn_select_memory(::jconf::inst()->GetMiningAlgo()), - cn_select_memory(::jconf::inst()->GetMiningAlgoRoot()) + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()), + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot()) ) / 1024u; const size_t halfHashMemSizeKB = hashMemSizeKB / 2u; diff --git a/xmrstak/backend/cpu/autoAdjustHwloc.hpp b/xmrstak/backend/cpu/autoAdjustHwloc.hpp index f110ee3ef..b1f391473 100644 --- a/xmrstak/backend/cpu/autoAdjustHwloc.hpp +++ b/xmrstak/backend/cpu/autoAdjustHwloc.hpp @@ -29,8 +29,8 @@ class autoAdjust autoAdjust() { hashMemSize = std::max( - cn_select_memory(::jconf::inst()->GetMiningAlgo()), - cn_select_memory(::jconf::inst()->GetMiningAlgoRoot()) + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()), + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot()) ); halfHashMemSize = hashMemSize / 2u; } diff --git a/xmrstak/backend/cpu/crypto/c_blake256.c b/xmrstak/backend/cpu/crypto/c_blake256.c index ff623ddb8..e5fadfe74 100644 --- a/xmrstak/backend/cpu/crypto/c_blake256.c +++ b/xmrstak/backend/cpu/crypto/c_blake256.c @@ -124,7 +124,7 @@ void blake224_init(state *S) { } // datalen = number of bits -void blake256_update(state *S, const uint8_t *data, uint64_t datalen) { +void blake256_update(state *S, const uint8_t *data, uint32_t datalen) { int left = S->buflen >> 3; int fill = 64 - left; @@ -155,7 +155,7 @@ void blake256_update(state *S, const uint8_t *data, uint64_t datalen) { } // datalen = number of bits -void blake224_update(state *S, const uint8_t *data, uint64_t datalen) { +void blake224_update(state *S, const uint8_t *data, uint32_t datalen) { blake256_update(S, data, datalen); } @@ -206,7 +206,7 @@ void blake224_final(state *S, uint8_t *digest) { } // inlen = number of bytes -void blake256_hash(uint8_t *out, const uint8_t *in, uint64_t inlen) { +void blake256_hash(uint8_t *out, const uint8_t *in, uint32_t inlen) { state S; blake256_init(&S); blake256_update(&S, in, inlen * 8); @@ -214,7 +214,7 @@ void blake256_hash(uint8_t *out, const uint8_t *in, uint64_t inlen) { } // inlen = number of bytes -void blake224_hash(uint8_t *out, const uint8_t *in, uint64_t inlen) { +void blake224_hash(uint8_t *out, const uint8_t *in, uint32_t inlen) { state S; blake224_init(&S); blake224_update(&S, in, inlen * 8); @@ -282,13 +282,13 @@ void hmac_blake224_init(hmac_state *S, const uint8_t *_key, uint64_t keylen) { } // datalen = number of bits -void hmac_blake256_update(hmac_state *S, const uint8_t *data, uint64_t datalen) { +void hmac_blake256_update(hmac_state *S, const uint8_t *data, uint32_t datalen) { // update the inner state blake256_update(&S->inner, data, datalen); } // datalen = number of bits -void hmac_blake224_update(hmac_state *S, const uint8_t *data, uint64_t datalen) { +void hmac_blake224_update(hmac_state *S, const uint8_t *data, uint32_t datalen) { // update the inner state blake224_update(&S->inner, data, datalen); } @@ -310,7 +310,7 @@ void hmac_blake224_final(hmac_state *S, uint8_t *digest) { } // keylen = number of bytes; inlen = number of bytes -void hmac_blake256_hash(uint8_t *out, const uint8_t *key, uint64_t keylen, const uint8_t *in, uint64_t inlen) { +void hmac_blake256_hash(uint8_t *out, const uint8_t *key, uint64_t keylen, const uint8_t *in, uint32_t inlen) { hmac_state S; hmac_blake256_init(&S, key, keylen); hmac_blake256_update(&S, in, inlen * 8); @@ -318,7 +318,7 @@ void hmac_blake256_hash(uint8_t *out, const uint8_t *key, uint64_t keylen, const } // keylen = number of bytes; inlen = number of bytes -void hmac_blake224_hash(uint8_t *out, const uint8_t *key, uint64_t keylen, const uint8_t *in, uint64_t inlen) { +void hmac_blake224_hash(uint8_t *out, const uint8_t *key, uint64_t keylen, const uint8_t *in, uint32_t inlen) { hmac_state S; hmac_blake224_init(&S, key, keylen); hmac_blake224_update(&S, in, inlen * 8); diff --git a/xmrstak/backend/cpu/crypto/c_blake256.h b/xmrstak/backend/cpu/crypto/c_blake256.h index b9c2aad0d..06c7917af 100644 --- a/xmrstak/backend/cpu/crypto/c_blake256.h +++ b/xmrstak/backend/cpu/crypto/c_blake256.h @@ -17,27 +17,27 @@ typedef struct { void blake256_init(state *); void blake224_init(state *); -void blake256_update(state *, const uint8_t *, uint64_t); -void blake224_update(state *, const uint8_t *, uint64_t); +void blake256_update(state *, const uint8_t *, uint32_t); +void blake224_update(state *, const uint8_t *, uint32_t); void blake256_final(state *, uint8_t *); void blake224_final(state *, uint8_t *); -void blake256_hash(uint8_t *, const uint8_t *, uint64_t); -void blake224_hash(uint8_t *, const uint8_t *, uint64_t); +void blake256_hash(uint8_t *, const uint8_t *, uint32_t); +void blake224_hash(uint8_t *, const uint8_t *, uint32_t); /* HMAC functions: */ void hmac_blake256_init(hmac_state *, const uint8_t *, uint64_t); void hmac_blake224_init(hmac_state *, const uint8_t *, uint64_t); -void hmac_blake256_update(hmac_state *, const uint8_t *, uint64_t); -void hmac_blake224_update(hmac_state *, const uint8_t *, uint64_t); +void hmac_blake256_update(hmac_state *, const uint8_t *, uint32_t); +void hmac_blake224_update(hmac_state *, const uint8_t *, uint32_t); void hmac_blake256_final(hmac_state *, uint8_t *); void hmac_blake224_final(hmac_state *, uint8_t *); -void hmac_blake256_hash(uint8_t *, const uint8_t *, uint64_t, const uint8_t *, uint64_t); -void hmac_blake224_hash(uint8_t *, const uint8_t *, uint64_t, const uint8_t *, uint64_t); +void hmac_blake256_hash(uint8_t *, const uint8_t *, uint64_t, const uint8_t *, uint32_t); +void hmac_blake224_hash(uint8_t *, const uint8_t *, uint64_t, const uint8_t *, uint32_t); #endif /* _BLAKE256_H_ */ diff --git a/xmrstak/backend/cpu/crypto/c_skein.h b/xmrstak/backend/cpu/crypto/c_skein.h index 6165a2ace..86dbc0802 100644 --- a/xmrstak/backend/cpu/crypto/c_skein.h +++ b/xmrstak/backend/cpu/crypto/c_skein.h @@ -37,7 +37,7 @@ typedef enum } SkeinHashReturn; -typedef size_t SkeinDataLength; /* bit count type */ +typedef uint32_t SkeinDataLength; /* bit count type */ typedef u08b_t SkeinBitSequence; /* bit stream type */ /* "all-in-one" call */ diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h index 6b555297d..7562de1bf 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h +++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h @@ -43,7 +43,7 @@ extern "C" { void keccak(const uint8_t *in, int inlen, uint8_t *md, int mdlen); void keccakf(uint64_t st[25], int rounds); - extern void(*const extra_hashes[4])(const void *, size_t, char *); + extern void(*const extra_hashes[4])(const void *, uint32_t, char *); } // This will shift and xor tmp1 into itself as 4 32-bit vals such as @@ -429,7 +429,7 @@ inline void cryptonight_monero_tweak(uint64_t* mem_out, __m128i 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; + uint8_t x = static_cast(vh >> 24); static const uint16_t table = 0x7531; const uint8_t index = (((x >> 3) & 6) | (x & 1)) << 1; vh ^= ((table >> index) & 0x3) << 28; @@ -444,7 +444,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c constexpr size_t ITERATIONS = cn_select_iter(); constexpr size_t MEM = cn_select_memory(); - if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) && 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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) { monero_const = *reinterpret_cast(reinterpret_cast(input) + 35); monero_const ^= *(reinterpret_cast(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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) + 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,15 +506,10 @@ 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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) ((uint64_t*)&l0[idx0 & MASK])[1] = ah0 ^ monero_const; else ((uint64_t*)&l0[idx0 & MASK])[1] = ah0; - - if (ALGO == cryptonight_ipbc) { - ((uint64_t*)&l0[idx0 & MASK])[1] ^= ((uint64_t*)&l0[idx0 & MASK])[0]; - } - ah0 ^= ch; idx0 = al0; @@ -549,7 +544,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto constexpr size_t ITERATIONS = cn_select_iter(); constexpr size_t MEM = cn_select_memory(); - if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43) { memset(output, 0, 64); return; @@ -559,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) { monero_const_0 = *reinterpret_cast(reinterpret_cast(input) + 35); monero_const_0 ^= *(reinterpret_cast(ctx[0]->hash_state) + 24); @@ -597,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) + 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)); @@ -615,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) + 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)); @@ -636,15 +631,11 @@ 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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) + 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; - if (ALGO == cryptonight_ipbc) { - ((uint64_t*)&l0[idx0 & MASK])[1] ^= ((uint64_t*)&l0[idx0 & MASK])[0]; - } - axh0 ^= ch; axl0 ^= cl; idx0 = axl0; @@ -671,15 +662,11 @@ 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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) + 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; - if (ALGO == cryptonight_ipbc) { - ((uint64_t*)&l1[idx1 & MASK])[1] ^= ((uint64_t*)&l1[idx1 & MASK])[0]; - } - axh1 ^= ch; axl1 ^= cl; idx1 = axl1; @@ -722,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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) \ + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) \ cryptonight_monero_tweak((uint64_t*)ptr, b); \ else \ _mm_store_si128(ptr, b);\ @@ -737,12 +724,10 @@ 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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) \ + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) \ _mm_store_si128(ptr, _mm_xor_si128(a, mc)); \ else \ _mm_store_si128(ptr, a);\ - if (ALGO == cryptonight_ipbc) \ - ((uint64_t*)&l[idx & MASK])[1] ^= ((uint64_t*)&l[idx & MASK])[0];\ a = _mm_xor_si128(a, b); \ idx = _mm_cvtsi128_si64(a); \ if(ALGO == cryptonight_heavy) \ @@ -766,7 +751,7 @@ void cryptonight_triple_hash(const void* input, size_t len, void* output, crypto constexpr size_t ITERATIONS = cn_select_iter(); constexpr size_t MEM = cn_select_memory(); - if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43) { memset(output, 0, 32 * 3); return; @@ -860,7 +845,7 @@ void cryptonight_quad_hash(const void* input, size_t len, void* output, cryptoni constexpr size_t ITERATIONS = cn_select_iter(); constexpr size_t MEM = cn_select_memory(); - if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43) { memset(output, 0, 32 * 4); return; @@ -969,7 +954,7 @@ void cryptonight_penta_hash(const void* input, size_t len, void* output, crypton constexpr size_t ITERATIONS = cn_select_iter(); constexpr size_t MEM = cn_select_memory(); - if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) && len < 43) + if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon) && len < 43) { memset(output, 0, 32 * 5); return; diff --git a/xmrstak/backend/cpu/crypto/cryptonight_common.cpp b/xmrstak/backend/cpu/crypto/cryptonight_common.cpp index ac696ddd3..ee3b66301 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight_common.cpp +++ b/xmrstak/backend/cpu/crypto/cryptonight_common.cpp @@ -56,23 +56,23 @@ extern "C" #include #endif // _WIN32 -void do_blake_hash(const void* input, size_t len, char* output) { +void do_blake_hash(const void* input, uint32_t len, char* output) { blake256_hash((uint8_t*)output, (const uint8_t*)input, len); } -void do_groestl_hash(const void* input, size_t len, char* output) { +void do_groestl_hash(const void* input, uint32_t len, char* output) { groestl((const uint8_t*)input, len * 8, (uint8_t*)output); } -void do_jh_hash(const void* input, size_t len, char* output) { +void do_jh_hash(const void* input, uint32_t len, char* output) { jh_hash(32 * 8, (const uint8_t*)input, 8 * len, (uint8_t*)output); } -void do_skein_hash(const void* input, size_t len, char* output) { +void do_skein_hash(const void* input, uint32_t len, char* output) { skein_hash(8 * 32, (const uint8_t*)input, 8 * len, (uint8_t*)output); } -void (* const extra_hashes[4])(const void *, size_t, char *) = {do_blake_hash, do_groestl_hash, do_jh_hash, do_skein_hash}; +void (* const extra_hashes[4])(const void *, uint32_t, char *) = {do_blake_hash, do_groestl_hash, do_jh_hash, do_skein_hash}; #ifdef _WIN32 #include "xmrstak/misc/uac.hpp" @@ -204,8 +204,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 = std::max( - cn_select_memory(::jconf::inst()->GetMiningAlgo()), - cn_select_memory(::jconf::inst()->GetMiningAlgoRoot()) + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()), + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot()) ); cryptonight_ctx* ptr = (cryptonight_ctx*)_mm_malloc(sizeof(cryptonight_ctx), 4096); @@ -283,8 +283,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 = std::max( - cn_select_memory(::jconf::inst()->GetMiningAlgo()), - cn_select_memory(::jconf::inst()->GetMiningAlgoRoot()) + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()), + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot()) ); if(ctx->ctx_info[0] != 0) diff --git a/xmrstak/backend/cpu/crypto/hash.h b/xmrstak/backend/cpu/crypto/hash.h index c12d355f4..2af330932 100644 --- a/xmrstak/backend/cpu/crypto/hash.h +++ b/xmrstak/backend/cpu/crypto/hash.h @@ -1,5 +1,7 @@ #pragma once +#include + typedef unsigned char BitSequence; -typedef unsigned long long DataLength; +typedef uint32_t DataLength; typedef enum {SUCCESS = 0, FAIL = 1, BAD_HASHLEN = 2} HashReturn; diff --git a/xmrstak/backend/cpu/hwlocMemory.cpp b/xmrstak/backend/cpu/hwlocMemory.cpp index 94d2b53f7..089570fc0 100644 --- a/xmrstak/backend/cpu/hwlocMemory.cpp +++ b/xmrstak/backend/cpu/hwlocMemory.cpp @@ -30,7 +30,7 @@ void bindMemoryToNUMANode( size_t puId ) depth = hwloc_get_type_depth(topology, HWLOC_OBJ_PU); - for( size_t i = 0; + for( uint32_t i = 0; i < hwloc_get_nbobjs_by_depth(topology, depth); i++ ) { diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp index 2046bced2..6937d9b8d 100644 --- a/xmrstak/backend/cpu/minethd.cpp +++ b/xmrstak/backend/cpu/minethd.cpp @@ -95,6 +95,7 @@ bool minethd::thd_setaffinity(std::thread::native_handle_type h, uint64_t cpu_id return pthread_setaffinity_np(h, sizeof(cpuset_t), &mn) == 0; #elif defined(__OpenBSD__) printer::inst()->print_msg(L0,"WARNING: thread pinning is not supported under OPENBSD."); + return true; #else cpu_set_t mn; CPU_ZERO(&mn); @@ -231,7 +232,7 @@ bool minethd::self_test() bool bResult = true; - if(::jconf::inst()->GetMiningAlgo() == cryptonight) + if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight) { unsigned char out[32 * MAX_N]; cn_hash_fun hashf; @@ -276,25 +277,15 @@ 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()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_lite) { } - else if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero) + else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_monero) { } - else if(::jconf::inst()->GetMiningAlgo() == cryptonight_aeon) + else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_aeon) { } - else if (::jconf::inst()->GetMiningAlgo() == cryptonight_ipbc) - { - unsigned char out[32 * MAX_N]; - cn_hash_fun hashf; - //cn_hash_fun_multi hashf_multi; - - hashf = func_selector(::jconf::inst()->HaveHardwareAes(), false, xmrstak_algo::cryptonight_ipbc); - hashf("This is a test PAAAAAAAAAAAAAAAAAAAAAAAAAAAAADDING", 50, out, ctx[0]); - bResult = memcmp(out, "\xa3\xed\x4b\x82\x52\x47\x95\x2d\x0e\x91\xd6\x46\xa3\x70\x9b\xa6\xe9\xf6\xab\x06\x39\x0b\x81\x81\x61\x1a\x77\xb6\x3c\x7d\x23\x88", 32) == 0; // \x8f\xd2\x87\x05\x1a\x60\x18\xcd\xf1\x11\x5a\xa0\xa6\xe2\x44\x86\x33\x97\xea\x31\x89\x59\x34\x60\x60\x75\xda\xb3\xca\x73\x6a\x24 - } for (int i = 0; i < MAX_N; i++) cryptonight_free_ctx(ctx[i]); @@ -382,9 +373,6 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr case cryptonight_aeon: algv = 4; break; - case cryptonight_ipbc: - algv = 5; - break; default: algv = 2; break; @@ -410,11 +398,7 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr cryptonight_hash, cryptonight_hash, cryptonight_hash, - cryptonight_hash, - cryptonight_hash, - cryptonight_hash, - cryptonight_hash, - cryptonight_hash + cryptonight_hash }; std::bitset<2> digit; @@ -441,7 +425,7 @@ void minethd::work_main() job_result result; // start with root algorithm and switch later if fork version is reached - auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot(); + auto miner_algo = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot(); cn_hash_fun hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo); ctx = minethd_alloc_ctx(); @@ -451,6 +435,7 @@ void minethd::work_main() result.iThreadId = iThreadNo; uint8_t version = 0; + size_t lastPoolId = 0; while (bQuit == 0) { @@ -478,16 +463,20 @@ void minethd::work_main() result.iNonce = *piNonce; uint8_t new_version = oWork.getVersion(); - - if (::jconf::inst()->GetMiningAlgo() == cryptonight_ipbc) new_version = oWork.bWorkBlob[1]; - - if(new_version != version) + if(new_version != version || oWork.iPoolId != lastPoolId) { - if(new_version >= ::jconf::inst()->GetMiningForkVersion()) + coinDescription coinDesc = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(oWork.iPoolId); + if(new_version >= coinDesc.GetMiningForkVersion()) { - miner_algo = ::jconf::inst()->GetMiningAlgo(); + miner_algo = coinDesc.GetMiningAlgo(); hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo); } + else + { + miner_algo = coinDesc.GetMiningAlgoRoot(); + hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo); + } + lastPoolId = oWork.iPoolId; version = new_version; } @@ -546,9 +535,6 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, case cryptonight_aeon: algv = 4; break; - case cryptonight_ipbc: - algv = 5; - break; default: algv = 2; break; @@ -638,24 +624,7 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, cryptonight_penta_hash, cryptonight_penta_hash, cryptonight_penta_hash, - cryptonight_penta_hash, - - cryptonight_double_hash, - cryptonight_double_hash, - cryptonight_double_hash, - cryptonight_double_hash, - cryptonight_triple_hash, - cryptonight_triple_hash, - cryptonight_triple_hash, - cryptonight_triple_hash, - cryptonight_quad_hash, - cryptonight_quad_hash, - cryptonight_quad_hash, - cryptonight_quad_hash, - cryptonight_penta_hash, - cryptonight_penta_hash, - cryptonight_penta_hash, - cryptonight_penta_hash + cryptonight_penta_hash }; std::bitset<2> digit; @@ -667,22 +636,22 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes, void minethd::double_work_main() { - multiway_work_main<2>(); + multiway_work_main<2u>(); } void minethd::triple_work_main() { - multiway_work_main<3>(); + multiway_work_main<3u>(); } void minethd::quad_work_main() { - multiway_work_main<4>(); + multiway_work_main<4u>(); } void minethd::penta_work_main() { - multiway_work_main<5>(); + multiway_work_main<5u>(); } template @@ -696,7 +665,7 @@ void minethd::prep_multiway_work(uint8_t *bWorkBlob, uint32_t **piNonce) } } -template +template void minethd::multiway_work_main() { if(affinity >= 0) //-1 means no affinity @@ -729,9 +698,10 @@ void minethd::multiway_work_main() globalStates::inst().iConsumeCnt++; // start with root algorithm and switch later if fork version is reached - auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot(); + auto miner_algo = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot(); cn_hash_fun_multi hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo); uint8_t version = 0; + size_t lastPoolId = 0; while (bQuit == 0) { @@ -758,14 +728,20 @@ void minethd::multiway_work_main() iNonce = *piNonce[0]; uint8_t new_version = oWork.getVersion(); - if (::jconf::inst()->GetMiningAlgo() == cryptonight_ipbc) new_version = oWork.bWorkBlob[1]; - if(new_version != version) + if(new_version != version || oWork.iPoolId != lastPoolId) { - if(new_version >= ::jconf::inst()->GetMiningForkVersion()) + coinDescription coinDesc = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(oWork.iPoolId); + if(new_version >= coinDesc.GetMiningForkVersion()) + { + miner_algo = coinDesc.GetMiningAlgo(); + hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo); + } + else { - miner_algo = ::jconf::inst()->GetMiningAlgo(); + miner_algo = coinDesc.GetMiningAlgoRoot(); hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo); } + lastPoolId = oWork.iPoolId; version = new_version; } diff --git a/xmrstak/backend/cpu/minethd.hpp b/xmrstak/backend/cpu/minethd.hpp index 59583b672..85a95d159 100644 --- a/xmrstak/backend/cpu/minethd.hpp +++ b/xmrstak/backend/cpu/minethd.hpp @@ -35,7 +35,7 @@ class minethd : public iBackend minethd(miner_work& pWork, size_t iNo, int iMultiway, bool no_prefetch, int64_t affinity); - template + template void multiway_work_main(); template diff --git a/xmrstak/backend/cryptonight.hpp b/xmrstak/backend/cryptonight.hpp index d31222675..8a8e25944 100644 --- a/xmrstak/backend/cryptonight.hpp +++ b/xmrstak/backend/cryptonight.hpp @@ -10,8 +10,7 @@ enum xmrstak_algo cryptonight_lite = 2, cryptonight_monero = 3, cryptonight_heavy = 4, - cryptonight_aeon = 5, - cryptonight_ipbc = 6 + cryptonight_aeon = 5 }; // define aeon settings @@ -27,10 +26,6 @@ constexpr size_t CRYPTONIGHT_HEAVY_MEMORY = 4 * 1024 * 1024; constexpr uint32_t CRYPTONIGHT_HEAVY_MASK = 0x3FFFF0; constexpr uint32_t CRYPTONIGHT_HEAVY_ITER = 0x40000; -constexpr size_t CRYPTONIGHT_IPBC_MEMORY = 1 * 1024 * 1024; -constexpr uint32_t CRYPTONIGHT_IPBC_MASK = CRYPTONIGHT_IPBC_MEMORY - 16; -constexpr uint32_t CRYPTONIGHT_IPBC_ITER = CRYPTONIGHT_IPBC_MEMORY / 4; - template inline constexpr size_t cn_select_memory() { return 0; } @@ -49,9 +44,6 @@ inline constexpr size_t cn_select_memory() { return CRYPTONIG template<> inline constexpr size_t cn_select_memory() { return CRYPTONIGHT_LITE_MEMORY; } -template<> -inline constexpr size_t cn_select_memory() { return CRYPTONIGHT_IPBC_MEMORY; } - inline size_t cn_select_memory(xmrstak_algo algo) { @@ -65,8 +57,6 @@ inline size_t cn_select_memory(xmrstak_algo algo) return CRYPTONIGHT_LITE_MEMORY; case cryptonight_heavy: return CRYPTONIGHT_HEAVY_MEMORY; - case cryptonight_ipbc: - return CRYPTONIGHT_IPBC_MEMORY; default: return 0; } @@ -90,10 +80,6 @@ inline constexpr uint32_t cn_select_mask() { return CRYPTONIG template<> inline constexpr uint32_t cn_select_mask() { return CRYPTONIGHT_LITE_MASK; } -template<> -inline constexpr uint32_t cn_select_mask() { return CRYPTONIGHT_IPBC_MASK; } - - inline size_t cn_select_mask(xmrstak_algo algo) { switch(algo) @@ -106,8 +92,6 @@ inline size_t cn_select_mask(xmrstak_algo algo) return CRYPTONIGHT_LITE_MASK; case cryptonight_heavy: return CRYPTONIGHT_HEAVY_MASK; - case cryptonight_ipbc: - return CRYPTONIGHT_IPBC_MASK; default: return 0; } @@ -131,10 +115,6 @@ inline constexpr uint32_t cn_select_iter() { return CRYPTONIG template<> inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_LITE_ITER; } -template<> -inline constexpr uint32_t cn_select_iter() { return CRYPTONIGHT_IPBC_ITER; } - - inline size_t cn_select_iter(xmrstak_algo algo) { switch(algo) @@ -147,8 +127,6 @@ inline size_t cn_select_iter(xmrstak_algo algo) return CRYPTONIGHT_LITE_ITER; case cryptonight_heavy: return CRYPTONIGHT_HEAVY_ITER; - case cryptonight_ipbc: - return CRYPTONIGHT_IPBC_ITER; default: return 0; } diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp index e66581173..92f5f789b 100644 --- a/xmrstak/backend/nvidia/minethd.cpp +++ b/xmrstak/backend/nvidia/minethd.cpp @@ -239,7 +239,7 @@ void minethd::work_main() cpu_ctx = cpu::minethd::minethd_alloc_ctx(); // start with root algorithm and switch later if fork version is reached - auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot(); + auto miner_algo = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot(); cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); uint32_t iNonce; @@ -247,6 +247,7 @@ void minethd::work_main() globalStates::inst().iConsumeCnt++; uint8_t version = 0; + size_t lastPoolId = 0; while (bQuit == 0) { @@ -264,15 +265,20 @@ void minethd::work_main() continue; } uint8_t new_version = oWork.getVersion(); - if (::jconf::inst()->GetMiningAlgo() == cryptonight_ipbc) new_version = oWork.bWorkBlob[1]; - - if(new_version != version) + if(new_version != version || oWork.iPoolId != lastPoolId) { - if(new_version >= ::jconf::inst()->GetMiningForkVersion()) + coinDescription coinDesc = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(oWork.iPoolId); + if(new_version >= coinDesc.GetMiningForkVersion()) + { + miner_algo = coinDesc.GetMiningAlgo(); + hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); + } + else { - miner_algo = ::jconf::inst()->GetMiningAlgo(); + miner_algo = coinDesc.GetMiningAlgoRoot(); hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); } + lastPoolId = oWork.iPoolId; version = new_version; } diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu index 6e4408cc1..43740d295 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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) + 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 || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) { const uint32_t table = 0x75310U; const uint32_t index = ((z >> 26) & 12) | ((z >> 23) & 2); @@ -304,22 +304,14 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti res = *( (uint64_t *) t2 ) >> ( sub & 1 ? 32 : 0 ); - if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) { - const uint32_t tweaked_res = tweak1_2[sub & 1] ^ res; - + 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 ); - - if (ALGO == cryptonight_ipbc && sub == 2) { - uint64_t* dst = ((uint64_t*)(long_state + j)); - uint64_t cur = loadGlobal64(dst); - uint64_t prev = loadGlobal64(dst - 1); - storeGlobal64(dst, cur ^ prev); - } a = ( sub & 1 ? yy[1] : yy[0] ) ^ res; idx0 = shuffle<4>(sPtr,sub, a, 0); @@ -500,8 +492,5 @@ void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t { cryptonight_core_gpu_hash(ctx, startNonce); } - else if (miner_algo == cryptonight_ipbc) - { - cryptonight_core_gpu_hash(ctx, startNonce); - } + } diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu index f016ef40b..f192f01dd 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu @@ -280,14 +280,14 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx) CUDA_CHECK(ctx->device_id, cudaDeviceSetCacheConfig(cudaFuncCachePreferL1)); size_t hashMemSize = std::max( - cn_select_memory(::jconf::inst()->GetMiningAlgo()), - cn_select_memory(::jconf::inst()->GetMiningAlgoRoot()) + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()), + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot()) ); 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()) + if(cryptonight_heavy == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()) { // extent ctx_b to hold the state of idx0 ctx_b_size += sizeof(uint32_t) * wsize; @@ -580,8 +580,8 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) ctx->free_device_memory = freeMemory; size_t hashMemSize = std::max( - cn_select_memory(::jconf::inst()->GetMiningAlgo()), - cn_select_memory(::jconf::inst()->GetMiningAlgoRoot()) + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()), + cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot()) ); #ifdef WIN32 @@ -612,7 +612,7 @@ 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()) + if(cryptonight_heavy == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()) perThread += 50 * 4; // state double buffer size_t max_intensity = limitedMemory / perThread; diff --git a/xmrstak/cli/cli-miner.cpp b/xmrstak/cli/cli-miner.cpp index c7debf8ff..e135678b4 100644 --- a/xmrstak/cli/cli-miner.cpp +++ b/xmrstak/cli/cli-miner.cpp @@ -80,6 +80,8 @@ void help() #ifndef CONF_NO_OPENCL cout<<" --noAMD disable the AMD miner backend"<=argc ) + { + printer::inst()->print_msg(L0, "No argument for parameter '--openCLVendor' given"); + win_exit(); + return 1; + } + std::string vendor(argv[i]); + params::inst().openCLVendor = vendor; + if(vendor != "AMD" && vendor != "NVIDIA") + { + printer::inst()->print_msg(L0, "'--openCLVendor' must be 'AMD' or 'NVIDIA'"); + win_exit(); + return 1; + } + } else if(opName.compare("--noAMDCache") == 0) { params::inst().AMDCache = false; diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp index 67fcf3739..fd59210de 100644 --- a/xmrstak/jconf.cpp +++ b/xmrstak/jconf.cpp @@ -86,33 +86,37 @@ configVal oConfigValues[] = { constexpr size_t iConfigCnt = (sizeof(oConfigValues)/sizeof(oConfigValues[0])); -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[] = { - { "aeon7", cryptonight_aeon, cryptonight_lite, 7u, "mine.aeon-pool.com:5555" }, - { "croat", cryptonight, cryptonight, 0u, nullptr }, - { "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_monero, cryptonight, 8u, nullptr }, - { "haven", cryptonight_heavy, cryptonight, 2u, nullptr }, - { "intense", cryptonight, cryptonight, 0u, nullptr }, - { "ipbc", cryptonight_ipbc, cryptonight, 1u, 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 } +xmrstak::coin_selection coins[] = { + // name, userpool, devpool, default_pool_suggestion + { "aeon7", {cryptonight_aeon, cryptonight_lite, 7u}, {cryptonight_aeon, cryptonight_lite, 7u}, nullptr }, + { "alloy", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "bitcoal", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "bbscoin", {cryptonight_monero, cryptonight, 3u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "croat", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "cryptonight", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "cryptonight_heavy", {cryptonight_heavy, cryptonight_heavy, 0u}, {cryptonight_heavy, cryptonight_heavy, 0u}, nullptr }, + { "cryptonight_lite", {cryptonight_aeon, cryptonight_lite, 255u}, {cryptonight_aeon, cryptonight_lite, 7u}, nullptr }, + { "cryptonight_lite_v7", {cryptonight_lite, cryptonight_aeon, 255u}, {cryptonight_aeon, cryptonight_lite, 7u}, nullptr }, + { "cryptonight_v7", {cryptonight_monero, cryptonight_monero, 0u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "dero", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "dinastycoin", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "edollar", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "electroneum", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "graft", {cryptonight_monero, cryptonight, 8u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "haven", {cryptonight_heavy, cryptonight, 2u}, {cryptonight_heavy, cryptonight_heavy, 0u}, nullptr }, + { "intense", {cryptonight_monero, cryptonight, 4u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "italocoin", {cryptonight_monero, cryptonight_monero, 0u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "karbo", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "leviarcoin", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "masari", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "monero7", {cryptonight_monero, cryptonight_monero, 0u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "stellite", {cryptonight_monero, cryptonight, 3u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, + { "sumokoin", {cryptonight_heavy, cryptonight_heavy, 0u}, {cryptonight_heavy, cryptonight_heavy, 0u}, nullptr }, + { "turtlecoin", {cryptonight_lite, cryptonight_aeon, 255u}, {cryptonight_aeon, cryptonight_lite, 7u}, nullptr }, + { "ultranote", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr }, }; -constexpr size_t coin_alogo_size = (sizeof(coin_algos)/sizeof(coin_algos[0])); +constexpr size_t coin_alogo_size = (sizeof(coins)/sizeof(coins[0])); inline bool checkType(Type have, Type want) { @@ -316,7 +320,7 @@ void jconf::GetAlgoList(std::string& list) for(size_t i=0; i < coin_alogo_size; i++) { list += "\t- "; - list += coin_algos[i].coin_name; + list += coins[i].coin_name; list += "\n"; } } @@ -334,7 +338,7 @@ bool jconf::IsOnAlgoList(std::string& needle) for(size_t i=0; i < coin_alogo_size; i++) { - if(needle == coin_algos[i].coin_name) + if(needle == coins[i].coin_name) return true; } return false; @@ -346,10 +350,10 @@ const char* jconf::GetDefaultPool(const char* needle) for(size_t i=0; i < coin_alogo_size; i++) { - if(strcmp(needle, coin_algos[i].coin_name) == 0) + if(strcmp(needle, coins[i].coin_name) == 0) { - if(coin_algos[i].default_pool != nullptr) - return coin_algos[i].default_pool; + if(coins[i].default_pool != nullptr) + return coins[i].default_pool; else return default_example; } @@ -633,16 +637,14 @@ bool jconf::parse_config(const char* sFilename, const char* sFilenamePools) return false; } - if(ctmp == coin_algos[i].coin_name) + if(ctmp == coins[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; + currentCoin = coins[i]; break; } } - if(mining_algo == invalid_algo) + if(currentCoin.GetDescription(1).GetMiningAlgo() == invalid_algo) { std::string cl; GetAlgoList(cl); diff --git a/xmrstak/jconf.hpp b/xmrstak/jconf.hpp index 76b93ddae..102b70f54 100644 --- a/xmrstak/jconf.hpp +++ b/xmrstak/jconf.hpp @@ -1,7 +1,7 @@ #pragma once -#include "xmrstak/backend/cryptonight.hpp" #include "xmrstak/misc/environment.hpp" +#include "xmrstak/misc/coinDescription.hpp" #include "params.hpp" #include @@ -48,12 +48,8 @@ class jconf bool TlsSecureAlgos(); - inline xmrstak_algo GetMiningAlgo() const { return mining_algo; } + inline xmrstak::coin_selection GetCurrentCoinSelection() const { return currentCoin; } - inline xmrstak_algo GetMiningAlgoRoot() const { return mining_algo_root; } - - inline uint8_t GetMiningForkVersion() const { return mining_fork_version; } - std::string GetMiningCoin(); static void GetAlgoList(std::string& list); @@ -94,7 +90,5 @@ class jconf opaque_private* prv; bool bHaveAes; - xmrstak_algo mining_algo; - xmrstak_algo mining_algo_root; - uint8_t mining_fork_version; + xmrstak::coin_selection currentCoin; }; diff --git a/xmrstak/misc/coinDescription.hpp b/xmrstak/misc/coinDescription.hpp new file mode 100644 index 000000000..73d9a9548 --- /dev/null +++ b/xmrstak/misc/coinDescription.hpp @@ -0,0 +1,60 @@ +#pragma once + +#include "xmrstak/backend/cryptonight.hpp" + +#include +#include + + +namespace xmrstak +{ + struct coinDescription + { + xmrstak_algo algo; + xmrstak_algo algo_root; + uint8_t fork_version; + + coinDescription() = default; + + coinDescription(const xmrstak_algo in_algo, xmrstak_algo in_algo_root, const uint8_t in_fork_version) : + algo(in_algo), algo_root(in_algo_root), fork_version(in_fork_version) + {} + + inline xmrstak_algo GetMiningAlgo() const { return algo; } + inline xmrstak_algo GetMiningAlgoRoot() const { return algo_root; } + inline uint8_t GetMiningForkVersion() const { return fork_version; } + }; + + struct coin_selection + { + const char* coin_name = nullptr; + /* [0] -> user pool + * [1] -> dev pool + */ + coinDescription pool_coin[2]; + const char* default_pool = nullptr; + + coin_selection() = default; + + coin_selection( + const char* in_coin_name, + const coinDescription user_coinDescription, + const coinDescription dev_coinDescription, + const char* in_default_pool + ) : + coin_name(in_coin_name), default_pool(in_default_pool) + { + pool_coin[0] = user_coinDescription; + pool_coin[1] = dev_coinDescription; + } + + /** get coin description for the pool + * + * @param poolId 0 select dev pool, else the user pool is selected + */ + inline coinDescription GetDescription(size_t poolId) const { + coinDescription tmp = (poolId == 0 ? pool_coin[1] : pool_coin[0]); + return tmp; + } + }; +} // namespace xmrstak diff --git a/xmrstak/misc/console.cpp b/xmrstak/misc/console.cpp index de5eed3db..d961b713b 100644 --- a/xmrstak/misc/console.cpp +++ b/xmrstak/misc/console.cpp @@ -225,7 +225,7 @@ void printer::print_str(const char* str) //Do a press any key for the windows folk. *insert any key joke here* #ifdef _WIN32 -void win_exit(size_t code) +void win_exit(int code) { size_t envSize = 0; getenv_s(&envSize, nullptr, 0, "XMRSTAK_NOWAIT"); @@ -238,7 +238,7 @@ void win_exit(size_t code) } #else -void win_exit(size_t code) +void win_exit(int code) { std::exit(code); } diff --git a/xmrstak/misc/console.hpp b/xmrstak/misc/console.hpp index cfbeddd01..671763105 100644 --- a/xmrstak/misc/console.hpp +++ b/xmrstak/misc/console.hpp @@ -49,4 +49,4 @@ class printer FILE* logfile; }; -void win_exit(size_t code = 1); +void win_exit(int code = 1); diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp index 61e826123..6e884ac3c 100644 --- a/xmrstak/misc/executor.cpp +++ b/xmrstak/misc/executor.cpp @@ -421,7 +421,9 @@ void executor::on_miner_result(size_t pool_id, job_result& oResult) { //Ignore errors silently if(pool->is_running() && pool->is_logged_in()) - pool->cmd_submit(oResult.sJobID, oResult.iNonce, oResult.bResult, backend_name, backend_hashcount, total_hashcount, jconf::inst()->GetMiningAlgo()); + pool->cmd_submit(oResult.sJobID, oResult.iNonce, oResult.bResult, backend_name, + backend_hashcount, total_hashcount, jconf::inst()->GetCurrentCoinSelection().GetDescription(0).GetMiningAlgo() + ); return; } @@ -432,7 +434,9 @@ 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, backend_name, backend_hashcount, total_hashcount, jconf::inst()->GetMiningAlgo()); + bool bResult = pool->cmd_submit(oResult.sJobID, oResult.iNonce, oResult.bResult, + backend_name, backend_hashcount, total_hashcount, jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() + ); size_t t_len = get_timestamp_ms() - t_start; if(t_len > 0xFFFF) @@ -548,7 +552,7 @@ 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); } - switch(jconf::inst()->GetMiningAlgo()) + switch(jconf::inst()->GetCurrentCoinSelection().GetDescription(0).GetMiningAlgo()) { case cryptonight_heavy: if(dev_tls) @@ -577,13 +581,6 @@ void executor::ex_main() else pools.emplace_front(0, "indeedminers.eu:1111", "", "", "", 0.0, true, false, "", true); break; - case cryptonight_ipbc: - if(dev_tls) - pools.emplace_front(0, "indeedminers.eu:7788", "", "", "", 0.0, true, false, "", true); - else - pools.emplace_front(0, "indeedminers.eu:7788", "", "", "", 0.0, true, false, "", true); - break; - default: break; diff --git a/xmrstak/misc/executor.hpp b/xmrstak/misc/executor.hpp index 9e145c3f5..be5ee6c2f 100644 --- a/xmrstak/misc/executor.hpp +++ b/xmrstak/misc/executor.hpp @@ -64,7 +64,7 @@ class executor inline bool is_dev_time() { //Add 2 seconds to compensate for connect - constexpr size_t dev_portion = double(iDevDonatePeriod) * fDevDonationLevel + 2; + constexpr size_t dev_portion = static_cast(double(iDevDonatePeriod) * fDevDonationLevel + 2.); if(dev_portion < 12) //No point in bothering with less than 10s return false; diff --git a/xmrstak/misc/telemetry.cpp b/xmrstak/misc/telemetry.cpp index 28fb522f9..a3a2a122f 100644 --- a/xmrstak/misc/telemetry.cpp +++ b/xmrstak/misc/telemetry.cpp @@ -89,8 +89,8 @@ double telemetry::calc_telemetry_data(size_t iLastMilisec, size_t iThread) return nan(""); double fHashes, fTime; - fHashes = iLastestHashCnt - iEarliestHashCnt; - fTime = iLastestStamp - iEarliestStamp; + fHashes = static_cast(iLastestHashCnt - iEarliestHashCnt); + fTime = static_cast(iLastestStamp - iEarliestStamp); fTime /= 1000.0; return fHashes / fTime; diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp index d71aeb1a6..74d1c261f 100644 --- a/xmrstak/net/jpsock.cpp +++ b/xmrstak/net/jpsock.cpp @@ -229,7 +229,7 @@ void jpsock::jpsock_thread() else disconnect_time = 0; - std::unique_lock(job_mutex); + std::unique_lock lck(job_mutex); memset(&oCurrentJob, 0, sizeof(oCurrentJob)); bRunning = false; } @@ -439,7 +439,7 @@ bool jpsock::process_pool_job(const opq_json_val* params) if(motd != nullptr && motd->IsString() && (motd->GetStringLength() & 0x01) == 0) { - std::unique_lock(motd_mutex); + std::unique_lock lck(motd_mutex); if(motd->GetStringLength() > 0) { pool_motd.resize(motd->GetStringLength()/2 + 1); @@ -454,7 +454,7 @@ bool jpsock::process_pool_job(const opq_json_val* params) executor::inst()->push_event(ex_event(oPoolJob, pool_id)); - std::unique_lock(job_mutex); + std::unique_lock lck(job_mutex); oCurrentJob = oPoolJob; return true; } @@ -679,13 +679,13 @@ bool jpsock::cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bRes void jpsock::save_nonce(uint32_t nonce) { - std::unique_lock(job_mutex); + std::unique_lock lck(job_mutex); oCurrentJob.iSavedNonce = nonce; } bool jpsock::get_current_job(pool_job& job) { - std::unique_lock(job_mutex); + std::unique_lock lck(job_mutex); if(oCurrentJob.iWorkLen == 0) return false; @@ -699,7 +699,7 @@ bool jpsock::get_pool_motd(std::string& strin) if(!ext_motd) return false; - std::unique_lock(motd_mutex); + std::unique_lock lck(motd_mutex); if(pool_motd.size() > 0) { strin.assign(pool_motd); diff --git a/xmrstak/net/socket.cpp b/xmrstak/net/socket.cpp index 7c58a8e06..9bc608fa7 100644 --- a/xmrstak/net/socket.cpp +++ b/xmrstak/net/socket.cpp @@ -156,7 +156,8 @@ int plain_socket::recv(char* buf, unsigned int len) bool plain_socket::send(const char* buf) { - int pos = 0, slen = strlen(buf); + size_t pos = 0; + size_t slen = strlen(buf); while (pos != slen) { int ret = ::send(hSocket, buf + pos, slen - pos, 0); diff --git a/xmrstak/params.hpp b/xmrstak/params.hpp index 0171f2884..6e676cf41 100644 --- a/xmrstak/params.hpp +++ b/xmrstak/params.hpp @@ -24,6 +24,8 @@ struct params bool AMDCache; bool useNVIDIA; bool useCPU; + // user selected OpenCL vendor + std::string openCLVendor; bool poolUseTls = false; std::string poolURL; @@ -60,6 +62,7 @@ struct params AMDCache(true), useNVIDIA(true), useCPU(true), + openCLVendor("AMD"), configFile("config.txt"), configFilePools("pools.txt"), configFileAMD("amd.txt"), diff --git a/xmrstak/pools.tpl b/xmrstak/pools.tpl index 847aca36d..5c1202704 100644 --- a/xmrstak/pools.tpl +++ b/xmrstak/pools.tpl @@ -18,22 +18,42 @@ R"===( POOLCONF], /* - * Currency to mine. Supported values: + * Currency to mine. (For IPBC download version 2.4.4) + * https://github.com/IndeedMiners/xmr-aeon-stak/releases/tag/2.4.4 + * Supported values: * * aeon7 (use this for Aeon's new PoW) + * alloy + * bbscoin (automatic switch with block version 3 to cryptonight_v7) + * bitcoal * croat - * cryptonight (try this if your coin is not listed) - * cryptonight_lite + * dero + * dinastycoin * edollar * electroneum * graft * haven * intense - * ipbc + * italocoin * karbo + * leviarcoin + * masari * monero7 (use this for Monero's new PoW) + * stellite * sumokoin + * turtlecoin + * ultranote + * + * Native algorithms which not depends on any block versions: * + * # 1MiB scratchpad memory + * cryptonight_lite + * cryptonight_lite_v7 + * # 2MiB scratchpad memory + * cryptonight + * cryptonight_v7 + * # 4MiB scratchpad memory + * cryptonight_heavy */ "currency" : "CURRENCY", diff --git a/xmrstak/version.cpp b/xmrstak/version.cpp index 939342b0a..bd043de31 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 dc15a069 +#define GIT_COMMIT_HASH a69a9c8d #endif #ifndef GIT_COMMIT_HASH @@ -19,7 +19,7 @@ #endif #define XMR_STAK_NAME "xmr-stak" -#define XMR_STAK_VERSION "2.4.4" +#define XMR_STAK_VERSION "2.4.5" #if defined(_WIN32) #define OS_TYPE "win"