diff options
Diffstat (limited to 'xmrstak/backend')
| -rw-r--r-- | xmrstak/backend/amd/amd_gpu/gpu.cpp | 54 | ||||
| -rw-r--r-- | xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl | 5 | ||||
| -rw-r--r-- | xmrstak/backend/amd/jconf.cpp | 10 | ||||
| -rw-r--r-- | xmrstak/backend/amd/minethd.cpp | 42 | ||||
| -rw-r--r-- | xmrstak/backend/cpu/autoAdjust.hpp | 14 | ||||
| -rw-r--r-- | xmrstak/backend/cpu/crypto/cryptonight_aesni.h | 30 | ||||
| -rw-r--r-- | xmrstak/backend/cpu/jconf.cpp | 10 | ||||
| -rw-r--r-- | xmrstak/backend/cpu/minethd.cpp | 120 | ||||
| -rw-r--r-- | xmrstak/backend/cpu/minethd.hpp | 2 | ||||
| -rw-r--r-- | xmrstak/backend/cryptonight.hpp | 26 | ||||
| -rw-r--r-- | xmrstak/backend/nvidia/jconf.cpp | 12 | ||||
| -rw-r--r-- | xmrstak/backend/nvidia/minethd.cpp | 47 | ||||
| -rw-r--r-- | xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp | 6 | ||||
| -rw-r--r-- | xmrstak/backend/nvidia/nvcc_code/cuda_core.cu | 24 | ||||
| -rw-r--r-- | xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu | 8 |
15 files changed, 214 insertions, 196 deletions
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, |
