diff --git a/.appveyor.yml b/.appveyor.yml index c336842d5..fee8f6524 100644 --- a/.appveyor.yml +++ b/.appveyor.yml @@ -12,7 +12,7 @@ install: - curl -sL https://github.com/fireice-uk/xmr-stak-dep/releases/download/v1/xmr-stak-dep.zip -o xmr-stak-dep.zip - 7z x xmr-stak-dep.zip -o"c:\xmr-stak-dep" -y > nul - appveyor DownloadFile https://developer.nvidia.com/compute/cuda/8.0/prod/local_installers/cuda_8.0.44_windows-exe -FileName cuda_8.0.44_windows.exe - - cuda_8.0.44_windows.exe -s compiler_8.0 cudart_8.0 + - cuda_8.0.44_windows.exe -s compiler_8.0 cudart_8.0 nvrtc_8.0 nvrtc_dev_8.0 - set PATH=%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v8.0\bin;%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v8.0\libnvvp;%PATH% - nvcc -V diff --git a/CMakeLists.txt b/CMakeLists.txt index 09ff7aef5..a5c06df8a 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -70,6 +70,42 @@ if(CUDA_ENABLE) find_package(CUDA 7.5) if(CUDA_FOUND) + # required for monero's cryptonight_r + # libcuda + find_library(CUDA_LIB + NAMES + libcuda + cuda + cuda.lib + HINTS + ${CUDA_TOOLKIT_ROOT_DIR} + ${LIBCUDA_LIBRARY_DIR} + ${CUDA_TOOLKIT_ROOT_DIR} + /usr + /usr/local/cuda + PATH_SUFFIXES + lib64 + lib/x64 + lib/Win32 + lib64/stubs) + + #nvrtc + find_library(CUDA_NVRTC_LIB + NAMES + libnvrtc + nvrtc + nvrtc.lib + HINTS + ${CUDA_TOOLKIT_ROOT_DIR} + ${LIBNVRTC_LIBRARY_DIR} + ${CUDA_TOOLKIT_ROOT_DIR} + /usr + /usr/local/cuda + PATH_SUFFIXES + lib64 + lib/x64 + lib/Win32) + list(APPEND BACKEND_TYPES "nvidia") option(XMR-STAK_LARGEGRID "Support large CUDA block count > 128" ON) if(XMR-STAK_LARGEGRID) @@ -200,16 +236,11 @@ if(CUDA_ENABLE) set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS}" "-D_MWAITXINTRIN_H_INCLUDED") endif() - if(CMAKE_CXX_COMPILER_ID MATCHES "MSVC" AND - (CUDA_VERSION VERSION_EQUAL 9.0 OR - CUDA_VERSION VERSION_EQUAL 9.1 OR - CUDA_VERSION VERSION_EQUAL 9.2 OR - CUDA_VERSION VERSION_EQUAL 10.0) - ) - # workaround find_package(CUDA) is using the wrong path to the CXX host compiler - # overwrite the CUDA host compiler variable with the used CXX MSVC - set(CUDA_HOST_COMPILER ${CMAKE_CXX_COMPILER} CACHE FILEPATH "Host side compiler used by NVCC" FORCE) - endif() + # workaround find_package(CUDA) is using the wrong path to the CXX host compiler + # overwrite the CUDA host compiler variable with the used CXX MSVC + # in linux where clang and gcc is installed it also helps to select the correct host compiler + set(CUDA_HOST_COMPILER ${CMAKE_CXX_COMPILER} CACHE FILEPATH "Host side compiler used by NVCC" FORCE) + else() message(FATAL_ERROR "selected CUDA compiler '${CUDA_COMPILER}' is not supported") endif() @@ -547,6 +578,8 @@ if(CUDA_FOUND) ${CUDASRCFILES} ) endif() + + set(CUDA_LIBRARIES ${CUDA_LIB} ${CUDA_NVRTC_LIB} ${CUDA_LIBRARIES}) target_link_libraries(xmrstak_cuda_backend ${CUDA_LIBRARIES}) target_link_libraries(xmrstak_cuda_backend xmr-stak-backend xmr-stak-asm) endif() diff --git a/xmrstak/backend/amd/OclCryptonightR_gen.cpp b/xmrstak/backend/amd/OclCryptonightR_gen.cpp new file mode 100644 index 000000000..4aabe51d0 --- /dev/null +++ b/xmrstak/backend/amd/OclCryptonightR_gen.cpp @@ -0,0 +1,354 @@ +#include +#include +#include +#include +#include + + +#include "xmrstak/backend/amd/OclCryptonightR_gen.hpp" +#include "xmrstak/backend/cpu/crypto/variant4_random_math.h" +#include "xmrstak/misc/console.hpp" +#include "xmrstak/cpputil/read_write_lock.h" + +#include +#include +#include + + +namespace xmrstak +{ +namespace amd +{ + +static std::string get_code(const V4_Instruction* code, int code_size) +{ + std::stringstream s; + + for (int i = 0; i < code_size; ++i) + { + const V4_Instruction inst = code[i]; + + const uint32_t a = inst.dst_index; + const uint32_t b = inst.src_index; + + switch (inst.opcode) + { + case MUL: + s << 'r' << a << "*=r" << b << ';'; + break; + + case ADD: + s << 'r' << a << "+=r" << b << '+' << inst.C << "U;"; + break; + + case SUB: + s << 'r' << a << "-=r" << b << ';'; + break; + + case ROR: + case ROL: + s << 'r' << a << "=rotate(r" << a << ((inst.opcode == ROR) ? ",ROT_BITS-r" : ",r") << b << ");"; + break; + + case XOR: + s << 'r' << a << "^=r" << b << ';'; + break; + } + + s << '\n'; + } + + return s.str(); +} + +struct CacheEntry +{ + CacheEntry(xmrstak_algo algo, uint64_t height, size_t deviceIdx, cl_program program) : + algo(algo), + height(height), + deviceIdx(deviceIdx), + program(program) + {} + + xmrstak_algo algo; + uint64_t height; + size_t deviceIdx; + cl_program program; +}; + +struct BackgroundTaskBase +{ + virtual ~BackgroundTaskBase() {} + virtual void exec() = 0; +}; + +template +struct BackgroundTask : public BackgroundTaskBase +{ + BackgroundTask(T&& func) : m_func(std::move(func)) {} + void exec() override { m_func(); } + + T m_func; +}; + +static ::cpputil::RWLock CryptonightR_cache_mutex; +static std::mutex CryptonightR_build_mutex; +static std::vector CryptonightR_cache; + +static std::mutex background_tasks_mutex; +static std::vector background_tasks; +static std::thread* background_thread = nullptr; + +static void background_thread_proc() +{ + std::vector tasks; + for (;;) { + tasks.clear(); + { + std::lock_guard g(background_tasks_mutex); + background_tasks.swap(tasks); + } + + for (BackgroundTaskBase* task : tasks) { + task->exec(); + delete task; + } + + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + } +} + +template +static void background_exec(T&& func) +{ + BackgroundTaskBase* task = new BackgroundTask(std::move(func)); + + std::lock_guard g(background_tasks_mutex); + background_tasks.push_back(task); + if (!background_thread) { + background_thread = new std::thread(background_thread_proc); + } +} + +static cl_program CryptonightR_build_program( + const GpuContext* ctx, + xmrstak_algo algo, + uint64_t height, + cl_kernel old_kernel, + std::string source_code, + std::string options) +{ + if(old_kernel) + clReleaseKernel(old_kernel); + + + std::vector old_programs; + old_programs.reserve(32); + { + CryptonightR_cache_mutex.WriteLock(); + + // Remove old programs from cache + for(size_t i = 0; i < CryptonightR_cache.size();) + { + const CacheEntry& entry = CryptonightR_cache[i]; + if ((entry.algo == algo) && (entry.height + 2 < height)) + { + printer::inst()->print_msg(LDEBUG, "CryptonightR: program for height %llu released (old program)", entry.height); + old_programs.push_back(entry.program); + CryptonightR_cache[i] = std::move(CryptonightR_cache.back()); + CryptonightR_cache.pop_back(); + } + else + { + ++i; + } + } + CryptonightR_cache_mutex.UnLock(); + } + + for(cl_program p : old_programs) { + clReleaseProgram(p); + } + + std::lock_guard g1(CryptonightR_build_mutex); + + cl_program program = nullptr; + { + CryptonightR_cache_mutex.ReadLock(); + + // Check if the cache already has this program (some other thread might have added it first) + for (const CacheEntry& entry : CryptonightR_cache) + { + if ((entry.algo == algo) && (entry.height == height) && (entry.deviceIdx == ctx->deviceIdx)) + { + program = entry.program; + break; + } + } + CryptonightR_cache_mutex.UnLock(); + } + + if (program) { + return program; + } + + cl_int ret; + const char* source = source_code.c_str(); + + program = clCreateProgramWithSource(ctx->opencl_ctx, 1, (const char**)&source, NULL, &ret); + if(ret != CL_SUCCESS) + { + printer::inst()->print_msg(L0,"Error %s when calling clCreateProgramWithSource on the OpenCL miner code", err_to_str(ret)); + return program; + } + + ret = clBuildProgram(program, 1, &ctx->DeviceID, options.c_str(), NULL, NULL); + if(ret != CL_SUCCESS) + { + size_t len; + printer::inst()->print_msg(L0,"Error %s when calling clBuildProgram.", err_to_str(ret)); + + if((ret = clGetProgramBuildInfo(program, ctx->DeviceID, CL_PROGRAM_BUILD_LOG, 0, NULL, &len)) != CL_SUCCESS) + { + printer::inst()->print_msg(L0,"Error %s when calling clGetProgramBuildInfo for length of build log output.", err_to_str(ret)); + return program; + } + + char* BuildLog = (char*)malloc(len + 1); + BuildLog[0] = '\0'; + + if((ret = clGetProgramBuildInfo(program, ctx->DeviceID, CL_PROGRAM_BUILD_LOG, len, BuildLog, NULL)) != CL_SUCCESS) + { + free(BuildLog); + printer::inst()->print_msg(L0,"Error %s when calling clGetProgramBuildInfo for build log.", err_to_str(ret)); + return program; + } + + printer::inst()->print_str("Build log:\n"); + std::cerr<DeviceID, CL_PROGRAM_BUILD_STATUS, sizeof(cl_build_status), &status, NULL)) != CL_SUCCESS) + { + printer::inst()->print_msg(L0,"Error %s when calling clGetProgramBuildInfo for status of build.", err_to_str(ret)); + return program; + } + std::this_thread::sleep_for(std::chrono::milliseconds(1000)); + } + while(status == CL_BUILD_IN_PROGRESS); + + + printer::inst()->print_msg(LDEBUG, "CryptonightR: program for height %llu compiled", height); + + CryptonightR_cache_mutex.WriteLock(); + CryptonightR_cache.emplace_back(algo, height, ctx->deviceIdx, program); + CryptonightR_cache_mutex.UnLock(); + return program; +} + +cl_program CryptonightR_get_program(GpuContext* ctx, xmrstak_algo algo, uint64_t height, bool background, cl_kernel old_kernel) +{ + if (background) { + background_exec([=](){ CryptonightR_get_program(ctx, algo, height, false, old_kernel); }); + return nullptr; + } + + const char* source_code_template = + #include "amd_gpu/opencl/wolf-aes.cl" + #include "amd_gpu/opencl/cryptonight_r.cl" + ; + const char include_name[] = "XMRSTAK_INCLUDE_RANDOM_MATH"; + const char* offset = strstr(source_code_template, include_name); + if (!offset) + { + printer::inst()->print_msg(LDEBUG, "CryptonightR_get_program: XMRSTAK_INCLUDE_RANDOM_MATH not found in cryptonight_r.cl", algo); + return nullptr; + } + + V4_Instruction code[256]; + int code_size; + switch (algo.Id()) + { + case cryptonight_r_wow: + code_size = v4_random_math_init(code, height); + break; + case cryptonight_r: + code_size = v4_random_math_init(code, height); + break; + default: + printer::inst()->print_msg(LDEBUG, "CryptonightR_get_program: invalid algo %d", algo); + return nullptr; + } + + std::string source_code(source_code_template, offset); + source_code.append(get_code(code, code_size)); + source_code.append(offset + sizeof(include_name) - 1); + + // scratchpad size for the selected mining algorithm + size_t hashMemSize = algo.Mem(); + int threadMemMask = algo.Mask(); + int hashIterations = algo.Iter(); + + size_t mem_chunk_exp = 1u << ctx->memChunk; + size_t strided_index = ctx->stridedIndex; + /* Adjust the config settings to a valid combination + * this is required if the dev pool is mining monero + * but the user tuned there settings for another currency + */ + if(algo == cryptonight_r || algo == cryptonight_r_wow) + { + if(ctx->memChunk < 2) + mem_chunk_exp = 1u << 2; + if(strided_index == 1) + strided_index = 0; + } + + // if intensity is a multiple of worksize than comp mode is not needed + int needCompMode = ctx->compMode && ctx->rawIntensity % ctx->workSize != 0 ? 1 : 0; + + std::string options; + options += " -DITERATIONS=" + std::to_string(hashIterations); + options += " -DMASK=" + std::to_string(threadMemMask) + "U"; + options += " -DWORKSIZE=" + std::to_string(ctx->workSize) + "U"; + options += " -DSTRIDED_INDEX=" + std::to_string(strided_index); + options += " -DMEM_CHUNK_EXPONENT=" + std::to_string(mem_chunk_exp) + "U"; + options += " -DCOMP_MODE=" + std::to_string(needCompMode); + options += " -DMEMORY=" + std::to_string(hashMemSize) + "LU"; + options += " -DALGO=" + std::to_string(algo.Id()); + options += " -DCN_UNROLL=" + std::to_string(ctx->unroll); + + if(algo == cryptonight_gpu) + options += " -cl-fp32-correctly-rounded-divide-sqrt"; + + + const char* source = source_code.c_str(); + + { + CryptonightR_cache_mutex.ReadLock(); + + // Check if the cache has this program + for (const CacheEntry& entry : CryptonightR_cache) + { + if ((entry.algo == algo) && (entry.height == height) && (entry.deviceIdx == ctx->deviceIdx)) + { + printer::inst()->print_msg(LDEBUG, "CryptonightR: program for height %llu found in cache", height); + auto result = entry.program; + CryptonightR_cache_mutex.UnLock(); + return result; + } + } + CryptonightR_cache_mutex.UnLock(); + + } + + return CryptonightR_build_program(ctx, algo, height, old_kernel, source, options); +} + +} // namespace amd +} // namespace xmrstak diff --git a/xmrstak/backend/amd/OclCryptonightR_gen.hpp b/xmrstak/backend/amd/OclCryptonightR_gen.hpp new file mode 100644 index 000000000..a69df9074 --- /dev/null +++ b/xmrstak/backend/amd/OclCryptonightR_gen.hpp @@ -0,0 +1,26 @@ +#pragma once + +#include "xmrstak/backend/cryptonight.hpp" + +#include +#include +#include + +#if defined(__APPLE__) +#include +#else +#include +#endif + +#include "xmrstak/backend/amd/amd_gpu/gpu.hpp" + +namespace xmrstak +{ +namespace amd +{ + +cl_program CryptonightR_get_program(GpuContext* ctx, const xmrstak_algo algo, + uint64_t height, bool background = false, cl_kernel old_kernel = nullptr); + +} // namespace amd +} // namespace xmrstak diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index 8713784c2..a2cbe8f54 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -19,6 +19,7 @@ #include "xmrstak/params.hpp" #include "xmrstak/version.hpp" #include "xmrstak/net/msgstruct.hpp" +#include "xmrstak/backend/amd/OclCryptonightR_gen.hpp" #include #include @@ -104,143 +105,6 @@ static inline long long unsigned int int_port(size_t i) #include "gpu.hpp" -const char* err_to_str(cl_int ret) -{ - switch(ret) - { - case CL_SUCCESS: - return "CL_SUCCESS"; - case CL_DEVICE_NOT_FOUND: - return "CL_DEVICE_NOT_FOUND"; - case CL_DEVICE_NOT_AVAILABLE: - return "CL_DEVICE_NOT_AVAILABLE"; - case CL_COMPILER_NOT_AVAILABLE: - return "CL_COMPILER_NOT_AVAILABLE"; - case CL_MEM_OBJECT_ALLOCATION_FAILURE: - return "CL_MEM_OBJECT_ALLOCATION_FAILURE"; - case CL_OUT_OF_RESOURCES: - return "CL_OUT_OF_RESOURCES"; - case CL_OUT_OF_HOST_MEMORY: - return "CL_OUT_OF_HOST_MEMORY"; - case CL_PROFILING_INFO_NOT_AVAILABLE: - return "CL_PROFILING_INFO_NOT_AVAILABLE"; - case CL_MEM_COPY_OVERLAP: - return "CL_MEM_COPY_OVERLAP"; - case CL_IMAGE_FORMAT_MISMATCH: - return "CL_IMAGE_FORMAT_MISMATCH"; - case CL_IMAGE_FORMAT_NOT_SUPPORTED: - return "CL_IMAGE_FORMAT_NOT_SUPPORTED"; - case CL_BUILD_PROGRAM_FAILURE: - return "CL_BUILD_PROGRAM_FAILURE"; - case CL_MAP_FAILURE: - return "CL_MAP_FAILURE"; - case CL_MISALIGNED_SUB_BUFFER_OFFSET: - return "CL_MISALIGNED_SUB_BUFFER_OFFSET"; - case CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST: - return "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST"; -#ifdef CL_VERSION_1_2 - case CL_COMPILE_PROGRAM_FAILURE: - return "CL_COMPILE_PROGRAM_FAILURE"; - case CL_LINKER_NOT_AVAILABLE: - return "CL_LINKER_NOT_AVAILABLE"; - case CL_LINK_PROGRAM_FAILURE: - return "CL_LINK_PROGRAM_FAILURE"; - case CL_DEVICE_PARTITION_FAILED: - return "CL_DEVICE_PARTITION_FAILED"; - case CL_KERNEL_ARG_INFO_NOT_AVAILABLE: - return "CL_KERNEL_ARG_INFO_NOT_AVAILABLE"; -#endif - case CL_INVALID_VALUE: - return "CL_INVALID_VALUE"; - case CL_INVALID_DEVICE_TYPE: - return "CL_INVALID_DEVICE_TYPE"; - case CL_INVALID_PLATFORM: - return "CL_INVALID_PLATFORM"; - case CL_INVALID_DEVICE: - return "CL_INVALID_DEVICE"; - case CL_INVALID_CONTEXT: - return "CL_INVALID_CONTEXT"; - case CL_INVALID_QUEUE_PROPERTIES: - return "CL_INVALID_QUEUE_PROPERTIES"; - case CL_INVALID_COMMAND_QUEUE: - return "CL_INVALID_COMMAND_QUEUE"; - case CL_INVALID_HOST_PTR: - return "CL_INVALID_HOST_PTR"; - case CL_INVALID_MEM_OBJECT: - return "CL_INVALID_MEM_OBJECT"; - case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR: - return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR"; - case CL_INVALID_IMAGE_SIZE: - return "CL_INVALID_IMAGE_SIZE"; - case CL_INVALID_SAMPLER: - return "CL_INVALID_SAMPLER"; - case CL_INVALID_BINARY: - return "CL_INVALID_BINARY"; - case CL_INVALID_BUILD_OPTIONS: - return "CL_INVALID_BUILD_OPTIONS"; - case CL_INVALID_PROGRAM: - return "CL_INVALID_PROGRAM"; - case CL_INVALID_PROGRAM_EXECUTABLE: - return "CL_INVALID_PROGRAM_EXECUTABLE"; - case CL_INVALID_KERNEL_NAME: - return "CL_INVALID_KERNEL_NAME"; - case CL_INVALID_KERNEL_DEFINITION: - return "CL_INVALID_KERNEL_DEFINITION"; - case CL_INVALID_KERNEL: - return "CL_INVALID_KERNEL"; - case CL_INVALID_ARG_INDEX: - return "CL_INVALID_ARG_INDEX"; - case CL_INVALID_ARG_VALUE: - return "CL_INVALID_ARG_VALUE"; - case CL_INVALID_ARG_SIZE: - return "CL_INVALID_ARG_SIZE"; - case CL_INVALID_KERNEL_ARGS: - return "CL_INVALID_KERNEL_ARGS"; - case CL_INVALID_WORK_DIMENSION: - return "CL_INVALID_WORK_DIMENSION"; - case CL_INVALID_WORK_GROUP_SIZE: - return "CL_INVALID_WORK_GROUP_SIZE"; - case CL_INVALID_WORK_ITEM_SIZE: - return "CL_INVALID_WORK_ITEM_SIZE"; - case CL_INVALID_GLOBAL_OFFSET: - return "CL_INVALID_GLOBAL_OFFSET"; - case CL_INVALID_EVENT_WAIT_LIST: - return "CL_INVALID_EVENT_WAIT_LIST"; - case CL_INVALID_EVENT: - return "CL_INVALID_EVENT"; - case CL_INVALID_OPERATION: - return "CL_INVALID_OPERATION"; - case CL_INVALID_GL_OBJECT: - return "CL_INVALID_GL_OBJECT"; - case CL_INVALID_BUFFER_SIZE: - return "CL_INVALID_BUFFER_SIZE"; - case CL_INVALID_MIP_LEVEL: - return "CL_INVALID_MIP_LEVEL"; - case CL_INVALID_GLOBAL_WORK_SIZE: - return "CL_INVALID_GLOBAL_WORK_SIZE"; - case CL_INVALID_PROPERTY: - return "CL_INVALID_PROPERTY"; -#ifdef CL_VERSION_1_2 - case CL_INVALID_IMAGE_DESCRIPTOR: - return "CL_INVALID_IMAGE_DESCRIPTOR"; - case CL_INVALID_COMPILER_OPTIONS: - return "CL_INVALID_COMPILER_OPTIONS"; - case CL_INVALID_LINKER_OPTIONS: - return "CL_INVALID_LINKER_OPTIONS"; - case CL_INVALID_DEVICE_PARTITION_COUNT: - return "CL_INVALID_DEVICE_PARTITION_COUNT"; -#endif -#if defined(CL_VERSION_2_0) && !defined(CONF_ENFORCE_OpenCL_1_2) - case CL_INVALID_PIPE_SIZE: - return "CL_INVALID_PIPE_SIZE"; - case CL_INVALID_DEVICE_QUEUE: - return "CL_INVALID_DEVICE_QUEUE"; -#endif - default: - return "UNKNOWN_ERROR"; - } -} - #if 0 void printer::inst()->print_msg(L1,const char* fmt, ...); void printer::inst()->print_str(const char* str); @@ -437,11 +301,20 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_ if(strided_index == 1) strided_index = 0; } + if(miner_algo == cryptonight_gpu) { strided_index = 0; } + if(miner_algo == cryptonight_r || miner_algo == cryptonight_r_wow) + { + if(ctx->memChunk < 2) + mem_chunk_exp = 1u << 2; + if(strided_index == 1) + strided_index = 0; + } + // if intensity is a multiple of worksize than comp mode is not needed int needCompMode = ctx->compMode && ctx->rawIntensity % ctx->workSize != 0 ? 1 : 0; @@ -853,8 +726,6 @@ int getAMDPlatformIdx() // Returns 0 on success, -1 on stupid params, -2 on OpenCL API error size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx) { - - cl_context opencl_ctx; cl_int ret; cl_uint entries; @@ -933,15 +804,6 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx) TempDeviceList[i] = DeviceIDList[ctx[i].deviceIdx]; } - opencl_ctx = clCreateContext(NULL, num_gpus, TempDeviceList, NULL, NULL, &ret); - if(ret != CL_SUCCESS) - { - printer::inst()->print_msg(L1,"Error %s when calling clCreateContext.", err_to_str(ret)); - return ERR_OCL_API; - } - - //char* source_code = LoadTextFile(sSourcePath); - const char *fastIntMathV2CL = #include "./opencl/fast_int_math_v2.cl" ; @@ -985,6 +847,20 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx) std::vector> interleaveData(num_gpus, nullptr); + std::vector context_vec(entries, nullptr); + for(int i = 0; i < num_gpus; ++i) + { + if(context_vec[ctx[i].deviceIdx] == nullptr) + { + context_vec[ctx[i].deviceIdx] = clCreateContext(NULL, 1, &(ctx[i].DeviceID), NULL, NULL, &ret); + if(ret != CL_SUCCESS) + { + printer::inst()->print_msg(L1,"Error %s when calling clCreateContext.", err_to_str(ret)); + return ERR_OCL_API; + } + } + } + for(int i = 0; i < num_gpus; ++i) { const size_t devIdx = ctx[i].deviceIdx; @@ -1003,8 +879,9 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx) ctx[i].interleaveData = interleaveData[devIdx]; ctx[i].interleaveData->adjustThreshold = static_cast(ctx[i].interleave)/100.0; ctx[i].interleaveData->startAdjustThreshold = ctx[i].interleaveData->adjustThreshold; + ctx[i].opencl_ctx = context_vec[ctx[i].deviceIdx]; - if((ret = InitOpenCLGpu(opencl_ctx, &ctx[i], source_code.c_str())) != ERR_SUCCESS) + if((ret = InitOpenCLGpu(ctx->opencl_ctx, &ctx[i], source_code.c_str())) != ERR_SUCCESS) { return ret; } @@ -1013,10 +890,10 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx) return ERR_SUCCESS; } -size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target, const xmrstak_algo& miner_algo) +size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target, const xmrstak_algo& miner_algo, uint64_t height) { - const auto & Kernels = ctx->Kernels[miner_algo.Id()]; + auto & Kernels = ctx->Kernels[miner_algo.Id()]; cl_int ret; @@ -1079,7 +956,41 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar } } - // CN1 Kernel + // CN1 Kernel + + if ((miner_algo == cryptonight_r) || (miner_algo == cryptonight_r_wow)) { + + // Get new kernel + cl_program program = xmrstak::amd::CryptonightR_get_program(ctx, miner_algo, height); + + if (program != ctx->ProgramCryptonightR) { + cl_int ret; + cl_kernel kernel = clCreateKernel(program, "cn1_cryptonight_r", &ret); + + cl_kernel old_kernel = nullptr; + if (ret != CL_SUCCESS) { + printer::inst()->print_msg(LDEBUG, "CryptonightR: clCreateKernel returned error %s", err_to_str(ret)); + } + else { + old_kernel = Kernels[1]; + Kernels[1] = kernel; + } + ctx->ProgramCryptonightR = program; + + uint32_t PRECOMPILATION_DEPTH = 4; + + // Precompile next program in background + xmrstak::amd::CryptonightR_get_program(ctx, miner_algo, height + 1, true, old_kernel); + for (int i = 2; i <= PRECOMPILATION_DEPTH; ++i) + xmrstak::amd::CryptonightR_get_program(ctx, miner_algo, height + i, true, nullptr); + + printer::inst()->print_msg(LDEBUG, "Thread #%zu updated CryptonightR", ctx->deviceIdx); + } + else + { + printer::inst()->print_msg(LDEBUG, "Thread #%zu found CryptonightR", ctx->deviceIdx); + } + } // Scratchpads if((ret = clSetKernelArg(Kernels[1], 0, sizeof(cl_mem), ctx->ExtraBuffers + 0)) != CL_SUCCESS) diff --git a/xmrstak/backend/amd/amd_gpu/gpu.hpp b/xmrstak/backend/amd/amd_gpu/gpu.hpp index d665dff1b..ae2b506db 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.hpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.hpp @@ -52,8 +52,10 @@ struct GpuContext cl_mem InputBuffer; cl_mem OutputBuffer; cl_mem ExtraBuffers[6]; + cl_context opencl_ctx = nullptr; std::map Program; std::map> Kernels; + cl_program ProgramCryptonightR = nullptr; size_t freeMem; size_t maxMemPerAlloc; int computeUnits; @@ -67,12 +69,152 @@ struct GpuContext }; +namespace +{ + const char* err_to_str(cl_int ret) + { + switch(ret) + { + case CL_SUCCESS: + return "CL_SUCCESS"; + case CL_DEVICE_NOT_FOUND: + return "CL_DEVICE_NOT_FOUND"; + case CL_DEVICE_NOT_AVAILABLE: + return "CL_DEVICE_NOT_AVAILABLE"; + case CL_COMPILER_NOT_AVAILABLE: + return "CL_COMPILER_NOT_AVAILABLE"; + case CL_MEM_OBJECT_ALLOCATION_FAILURE: + return "CL_MEM_OBJECT_ALLOCATION_FAILURE"; + case CL_OUT_OF_RESOURCES: + return "CL_OUT_OF_RESOURCES"; + case CL_OUT_OF_HOST_MEMORY: + return "CL_OUT_OF_HOST_MEMORY"; + case CL_PROFILING_INFO_NOT_AVAILABLE: + return "CL_PROFILING_INFO_NOT_AVAILABLE"; + case CL_MEM_COPY_OVERLAP: + return "CL_MEM_COPY_OVERLAP"; + case CL_IMAGE_FORMAT_MISMATCH: + return "CL_IMAGE_FORMAT_MISMATCH"; + case CL_IMAGE_FORMAT_NOT_SUPPORTED: + return "CL_IMAGE_FORMAT_NOT_SUPPORTED"; + case CL_BUILD_PROGRAM_FAILURE: + return "CL_BUILD_PROGRAM_FAILURE"; + case CL_MAP_FAILURE: + return "CL_MAP_FAILURE"; + case CL_MISALIGNED_SUB_BUFFER_OFFSET: + return "CL_MISALIGNED_SUB_BUFFER_OFFSET"; + case CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST: + return "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST"; + #ifdef CL_VERSION_1_2 + case CL_COMPILE_PROGRAM_FAILURE: + return "CL_COMPILE_PROGRAM_FAILURE"; + case CL_LINKER_NOT_AVAILABLE: + return "CL_LINKER_NOT_AVAILABLE"; + case CL_LINK_PROGRAM_FAILURE: + return "CL_LINK_PROGRAM_FAILURE"; + case CL_DEVICE_PARTITION_FAILED: + return "CL_DEVICE_PARTITION_FAILED"; + case CL_KERNEL_ARG_INFO_NOT_AVAILABLE: + return "CL_KERNEL_ARG_INFO_NOT_AVAILABLE"; + #endif + case CL_INVALID_VALUE: + return "CL_INVALID_VALUE"; + case CL_INVALID_DEVICE_TYPE: + return "CL_INVALID_DEVICE_TYPE"; + case CL_INVALID_PLATFORM: + return "CL_INVALID_PLATFORM"; + case CL_INVALID_DEVICE: + return "CL_INVALID_DEVICE"; + case CL_INVALID_CONTEXT: + return "CL_INVALID_CONTEXT"; + case CL_INVALID_QUEUE_PROPERTIES: + return "CL_INVALID_QUEUE_PROPERTIES"; + case CL_INVALID_COMMAND_QUEUE: + return "CL_INVALID_COMMAND_QUEUE"; + case CL_INVALID_HOST_PTR: + return "CL_INVALID_HOST_PTR"; + case CL_INVALID_MEM_OBJECT: + return "CL_INVALID_MEM_OBJECT"; + case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR: + return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR"; + case CL_INVALID_IMAGE_SIZE: + return "CL_INVALID_IMAGE_SIZE"; + case CL_INVALID_SAMPLER: + return "CL_INVALID_SAMPLER"; + case CL_INVALID_BINARY: + return "CL_INVALID_BINARY"; + case CL_INVALID_BUILD_OPTIONS: + return "CL_INVALID_BUILD_OPTIONS"; + case CL_INVALID_PROGRAM: + return "CL_INVALID_PROGRAM"; + case CL_INVALID_PROGRAM_EXECUTABLE: + return "CL_INVALID_PROGRAM_EXECUTABLE"; + case CL_INVALID_KERNEL_NAME: + return "CL_INVALID_KERNEL_NAME"; + case CL_INVALID_KERNEL_DEFINITION: + return "CL_INVALID_KERNEL_DEFINITION"; + case CL_INVALID_KERNEL: + return "CL_INVALID_KERNEL"; + case CL_INVALID_ARG_INDEX: + return "CL_INVALID_ARG_INDEX"; + case CL_INVALID_ARG_VALUE: + return "CL_INVALID_ARG_VALUE"; + case CL_INVALID_ARG_SIZE: + return "CL_INVALID_ARG_SIZE"; + case CL_INVALID_KERNEL_ARGS: + return "CL_INVALID_KERNEL_ARGS"; + case CL_INVALID_WORK_DIMENSION: + return "CL_INVALID_WORK_DIMENSION"; + case CL_INVALID_WORK_GROUP_SIZE: + return "CL_INVALID_WORK_GROUP_SIZE"; + case CL_INVALID_WORK_ITEM_SIZE: + return "CL_INVALID_WORK_ITEM_SIZE"; + case CL_INVALID_GLOBAL_OFFSET: + return "CL_INVALID_GLOBAL_OFFSET"; + case CL_INVALID_EVENT_WAIT_LIST: + return "CL_INVALID_EVENT_WAIT_LIST"; + case CL_INVALID_EVENT: + return "CL_INVALID_EVENT"; + case CL_INVALID_OPERATION: + return "CL_INVALID_OPERATION"; + case CL_INVALID_GL_OBJECT: + return "CL_INVALID_GL_OBJECT"; + case CL_INVALID_BUFFER_SIZE: + return "CL_INVALID_BUFFER_SIZE"; + case CL_INVALID_MIP_LEVEL: + return "CL_INVALID_MIP_LEVEL"; + case CL_INVALID_GLOBAL_WORK_SIZE: + return "CL_INVALID_GLOBAL_WORK_SIZE"; + case CL_INVALID_PROPERTY: + return "CL_INVALID_PROPERTY"; + #ifdef CL_VERSION_1_2 + case CL_INVALID_IMAGE_DESCRIPTOR: + return "CL_INVALID_IMAGE_DESCRIPTOR"; + case CL_INVALID_COMPILER_OPTIONS: + return "CL_INVALID_COMPILER_OPTIONS"; + case CL_INVALID_LINKER_OPTIONS: + return "CL_INVALID_LINKER_OPTIONS"; + case CL_INVALID_DEVICE_PARTITION_COUNT: + return "CL_INVALID_DEVICE_PARTITION_COUNT"; + #endif + #if defined(CL_VERSION_2_0) && !defined(CONF_ENFORCE_OpenCL_1_2) + case CL_INVALID_PIPE_SIZE: + return "CL_INVALID_PIPE_SIZE"; + case CL_INVALID_DEVICE_QUEUE: + return "CL_INVALID_DEVICE_QUEUE"; + #endif + default: + return "UNKNOWN_ERROR"; + } + } +} + uint32_t getNumPlatforms(); int getAMDPlatformIdx(); std::vector getAMDDevices(int index); size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx); -size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target, const xmrstak_algo& miner_algo); +size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target, const xmrstak_algo& miner_algo, uint64_t height); size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, const xmrstak_algo& miner_algo); uint64_t interleaveAdjustDelay(GpuContext* ctx, const bool enableAutoAdjustment = true); uint64_t updateTimings(GpuContext* ctx, const uint64_t t); diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight_r.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight_r.cl new file mode 100644 index 000000000..9edb774ad --- /dev/null +++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight_r.cl @@ -0,0 +1,220 @@ +R"===( +/* + * This program is free software: you can redistribute it and/or modify + * it under the terms of the GNU General Public License as published by + * the Free Software Foundation, either version 3 of the License, or + * any later version. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with this program. If not, see . + * + */ + +#define cryptonight_r_wow 15 +#define cryptonight_r 16 + +#define MEM_CHUNK (1 << MEM_CHUNK_EXPONENT) + +#if(STRIDED_INDEX==0) +# define IDX(x) (x) +#elif(STRIDED_INDEX==1) +# define IDX(x) (mul24(((uint)(x)), Threads)) +#elif(STRIDED_INDEX==2) +# define IDX(x) (((x) % MEM_CHUNK) + ((x) / MEM_CHUNK) * WORKSIZE * MEM_CHUNK) +#elif(STRIDED_INDEX==3) +# define IDX(x) ((x) * WORKSIZE) +#endif + +// __NV_CL_C_VERSION checks if NVIDIA opencl is used +#if(ALGO == cryptonight_monero_v8 && defined(__NV_CL_C_VERSION)) +# define SCRATCHPAD_CHUNK(N) (*(__local uint4*)((__local uchar*)(scratchpad_line) + (idx1 ^ (N << 4)))) +# define SCRATCHPAD_CHUNK_GLOBAL (*((__global uint16*)(Scratchpad + (IDX((idx0 & 0x1FFFC0U) >> 4))))) +#else +# define SCRATCHPAD_CHUNK(N) (Scratchpad[IDX(((idx) >> 4) ^ N)]) +#endif + +__attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) +__kernel void cn1_cryptonight_r(__global uint4 *Scratchpad, __global ulong *states, uint Threads) +{ + ulong a[2], b[4]; + __local uint AES0[256], AES1[256], AES2[256], AES3[256]; + +#ifdef __NV_CL_C_VERSION + __local uint16 scratchpad_line_buf[WORKSIZE]; + __local uint16* scratchpad_line = scratchpad_line_buf + get_local_id(0); +#endif + + const ulong gIdx = get_global_id(0) - get_global_offset(0); + + for(int i = get_local_id(0); i < 256; i += WORKSIZE) + { + const uint tmp = AES0_C[i]; + AES0[i] = tmp; + AES1[i] = rotate(tmp, 8U); + AES2[i] = rotate(tmp, 16U); + AES3[i] = rotate(tmp, 24U); + } + + barrier(CLK_LOCAL_MEM_FENCE); + +# if (COMP_MODE == 1) + // do not use early return here + if (gIdx < Threads) +# endif + { + states += 25 * gIdx; + +#if(STRIDED_INDEX==0) + Scratchpad += gIdx * (MEMORY >> 4); +#elif(STRIDED_INDEX==1) + Scratchpad += gIdx; +#elif(STRIDED_INDEX==2) + Scratchpad += get_group_id(0) * (MEMORY >> 4) * WORKSIZE + MEM_CHUNK * get_local_id(0); +#elif(STRIDED_INDEX==3) + Scratchpad += (gIdx / WORKSIZE) * (MEMORY >> 4) * WORKSIZE + (gIdx % WORKSIZE); +#endif + + a[0] = states[0] ^ states[4]; + a[1] = states[1] ^ states[5]; + + b[0] = states[2] ^ states[6]; + b[1] = states[3] ^ states[7]; + b[2] = states[8] ^ states[10]; + b[3] = states[9] ^ states[11]; + } + + ulong2 bx0 = ((ulong2 *)b)[0]; + ulong2 bx1 = ((ulong2 *)b)[1]; + + mem_fence(CLK_LOCAL_MEM_FENCE); + +# if (COMP_MODE == 1) + // do not use early return here + if (gIdx < Threads) +# endif + { + + uint r0 = as_uint2(states[12]).s0; + uint r1 = as_uint2(states[12]).s1; + uint r2 = as_uint2(states[13]).s0; + uint r3 = as_uint2(states[13]).s1; + + #pragma unroll CN_UNROLL + for(int i = 0; i < ITERATIONS; ++i) + { +# ifdef __NV_CL_C_VERSION + uint idx = a[0] & 0x1FFFC0; + uint idx1 = a[0] & 0x30; + + *scratchpad_line = *(__global uint16*)((__global uchar*)(Scratchpad) + idx); +# else + uint idx = a[0] & MASK; +# endif + +#if(ALGO == cryptonight_monero_v8 && defined(__NV_CL_C_VERSION)) + *scratchpad_line = SCRATCHPAD_CHUNK_GLOBAL; +#endif + uint4 c = SCRATCHPAD_CHUNK(0); + c = AES_Round(AES0, AES1, AES2, AES3, c, ((uint4 *)a)[0]); + + { + const ulong2 chunk1 = as_ulong2(SCRATCHPAD_CHUNK(1)); + const ulong2 chunk2 = as_ulong2(SCRATCHPAD_CHUNK(2)); + const ulong2 chunk3 = as_ulong2(SCRATCHPAD_CHUNK(3)); + +#if (ALGO == cryptonight_r) + c ^= as_uint4(chunk1) ^ as_uint4(chunk2) ^ as_uint4(chunk3); +#endif + + SCRATCHPAD_CHUNK(1) = as_uint4(chunk3 + bx1); + SCRATCHPAD_CHUNK(2) = as_uint4(chunk1 + bx0); + SCRATCHPAD_CHUNK(3) = as_uint4(chunk2 + ((ulong2 *)a)[0]); + } + + SCRATCHPAD_CHUNK(0) = as_uint4(bx0) ^ c; + +# ifdef __NV_CL_C_VERSION + *(__global uint16*)((__global uchar*)(Scratchpad) + idx) = *scratchpad_line; + + idx = as_ulong2(c).s0 & 0x1FFFC0; + idx1 = as_ulong2(c).s0 & 0x30; + + *scratchpad_line = *(__global uint16*)((__global uchar*)(Scratchpad) + idx); +# else + idx = as_ulong2(c).s0 & MASK; +# endif + + uint4 tmp = SCRATCHPAD_CHUNK(0); + + tmp.s0 ^= r0 + r1; + tmp.s1 ^= r2 + r3; + const uint r4 = as_uint2(a[0]).s0; + const uint r5 = as_uint2(a[1]).s0; + const uint r6 = as_uint4(bx0).s0; + const uint r7 = as_uint4(bx1).s0; +#if (ALGO == cryptonight_r) + const uint r8 = as_uint4(bx1).s2; +#endif +#define ROT_BITS 32 + + XMRSTAK_INCLUDE_RANDOM_MATH + +#if (ALGO == cryptonight_r) + + const uint2 al = (uint2)(as_uint2(a[0]).s0 ^ r2, as_uint2(a[0]).s1 ^ r3); + const uint2 ah = (uint2)(as_uint2(a[1]).s0 ^ r0, as_uint2(a[1]).s1 ^ r1); +#endif + + ulong2 t; + t.s0 = mul_hi(as_ulong2(c).s0, as_ulong2(tmp).s0); + t.s1 = as_ulong2(c).s0 * as_ulong2(tmp).s0; + { + const ulong2 chunk1 = as_ulong2(SCRATCHPAD_CHUNK(1)) +#if (ALGO == cryptonight_r_wow) + ^ t +#endif + ; + const ulong2 chunk2 = as_ulong2(SCRATCHPAD_CHUNK(2)); +#if (ALGO == cryptonight_r_wow) + t ^= chunk2; +#endif + const ulong2 chunk3 = as_ulong2(SCRATCHPAD_CHUNK(3)); + +#if (ALGO == cryptonight_r) + c ^= as_uint4(chunk1) ^ as_uint4(chunk2) ^ as_uint4(chunk3); +#endif + + SCRATCHPAD_CHUNK(1) = as_uint4(chunk3 + bx1); + SCRATCHPAD_CHUNK(2) = as_uint4(chunk1 + bx0); + SCRATCHPAD_CHUNK(3) = as_uint4(chunk2 + ((ulong2 *)a)[0]); + } + +#if (ALGO == cryptonight_r) + a[1] = as_ulong(ah) + t.s1; + a[0] = as_ulong(al) + t.s0; +#else + a[1] += t.s1; + a[0] += t.s0; +#endif + + SCRATCHPAD_CHUNK(0) = ((uint4 *)a)[0]; + +# ifdef __NV_CL_C_VERSION + *(__global uint16*)((__global uchar*)(Scratchpad) + idx) = *scratchpad_line; +# endif + + ((uint4 *)a)[0] ^= tmp; + bx1 = bx0; + bx0 = as_ulong2(c); + } + +# undef SCRATCHPAD_CHUNK + } + mem_fence(CLK_GLOBAL_MEM_FENCE); +} +)===" diff --git a/xmrstak/backend/amd/autoAdjust.hpp b/xmrstak/backend/amd/autoAdjust.hpp index 48f4ca49a..ea688e053 100644 --- a/xmrstak/backend/amd/autoAdjust.hpp +++ b/xmrstak/backend/amd/autoAdjust.hpp @@ -139,6 +139,10 @@ class autoAdjust // true for cryptonight_gpu as main user pool algorithm bool useCryptonight_gpu = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_gpu; + bool useCryptonight_r = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_r; + + bool useCryptonight_r_wow = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_r_wow; + // set strided index to default ctx.stridedIndex = 1; @@ -147,7 +151,7 @@ class autoAdjust ctx.stridedIndex = 0; // use chunked (4x16byte) scratchpad for all backends. Default `mem_chunk` is `2` - if(useCryptonight_v8) + if(useCryptonight_v8 || useCryptonight_r || useCryptonight_r_wow) ctx.stridedIndex = 2; else if(useCryptonight_heavy) ctx.stridedIndex = 3; diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp index 1c9eb6279..eb0009413 100644 --- a/xmrstak/backend/amd/minethd.cpp +++ b/xmrstak/backend/amd/minethd.cpp @@ -183,7 +183,11 @@ void minethd::work_main() } // start with root algorithm and switch later if fork version is reached 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); + + cpu::minethd::cn_on_new_job set_job; + + cn_hash_fun hash_fun; + cpu::minethd::func_multi_selector<1>(hash_fun, set_job, ::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); uint8_t version = 0; size_t lastPoolId = 0; @@ -224,23 +228,26 @@ void minethd::work_main() if(new_version >= coinDesc.GetMiningForkVersion()) { miner_algo = coinDesc.GetMiningAlgo(); - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); + cpu::minethd::func_multi_selector<1>(hash_fun, set_job, ::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); } else { miner_algo = coinDesc.GetMiningAlgoRoot(); - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); + cpu::minethd::func_multi_selector<1>(hash_fun, set_job, ::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); } lastPoolId = oWork.iPoolId; version = new_version; } + if(set_job != nullptr) + set_job(oWork, &cpu_ctx); + size_t round_ctr = 0; assert(sizeof(job_result::sJobID) == sizeof(pool_job::sJobID)); uint64_t target = oWork.iTarget; - XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target, miner_algo); + XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target, miner_algo, cpu_ctx->cn_r_ctx.height); if(oWork.bNiceHash) pGpuCtx->Nonce = *(uint32_t*)(oWork.bWorkBlob + 39); @@ -327,7 +334,7 @@ void minethd::work_main() ); } // update gpu with new intensity - XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target, miner_algo); + XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target, miner_algo, cpu_ctx->cn_r_ctx.height); } // use 3 rounds to warm up with the new intensity else if(cntTestRounds == autoTune + 3) diff --git a/xmrstak/backend/cpu/crypto/cryptonight.h b/xmrstak/backend/cpu/crypto/cryptonight.h index 5c9a73332..a7c77cdac 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight.h +++ b/xmrstak/backend/cpu/crypto/cryptonight.h @@ -1,29 +1,31 @@ -#ifndef __CRYPTONIGHT_H_INCLUDED -#define __CRYPTONIGHT_H_INCLUDED - -#ifdef __cplusplus -extern "C" { -#endif - +#pragma once #include #include -typedef struct { +#include "variant4_random_math.h" + +struct extra_ctx_r +{ + uint64_t height = 0; + // the buffer must be able to hold NUM_INSTRUCTIONS_MAX and a termination instruction + V4_Instruction code[NUM_INSTRUCTIONS_MAX + 1]; +}; + +struct cryptonight_ctx +{ uint8_t hash_state[224]; // Need only 200, explicit align uint8_t* long_state; uint8_t ctx_info[24]; //Use some of the extra memory for flags -} cryptonight_ctx; + extra_ctx_r cn_r_ctx; +}; -typedef struct { +struct alloc_msg +{ const char* warning; -} alloc_msg; +}; 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); void cryptonight_free_ctx(cryptonight_ctx* ctx); -#ifdef __cplusplus -} -#endif -#endif diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h index dc378e88a..43f719873 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h +++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h @@ -17,6 +17,7 @@ #include "cryptonight.h" #include "xmrstak/backend/cryptonight.hpp" +#include "../../miner_work.hpp" #include "cn_gpu.hpp" #include #include @@ -585,7 +586,7 @@ inline void set_float_rounding_mode() #endif } -inline void set_float_rounding_mode_conceal() +inline void set_float_rounding_mode_nearest() { #ifdef _MSC_VER _control87(RC_NEAR, MCW_RC); @@ -615,9 +616,9 @@ inline void cryptonight_conceal_tweak(__m128i& cx, __m128& conc_var) cx = _mm_xor_si128(cx, _mm_cvttps_epi32(nc)); } -#define CN_MONERO_V8_SHUFFLE_0(n, l0, idx0, ax0, bx0, bx1) \ +#define CN_MONERO_V8_SHUFFLE_0(n, l0, idx0, ax0, bx0, bx1, cx) \ /* Shuffle the other 3x16 byte chunks in the current 64-byte cache line */ \ - if(ALGO == cryptonight_monero_v8) \ + if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_r || ALGO == cryptonight_r_wow) \ { \ const uint64_t idx1 = idx0 & MASK; \ const __m128i chunk1 = _mm_load_si128((__m128i *)&l0[idx1 ^ 0x10]); \ @@ -626,11 +627,13 @@ inline void cryptonight_conceal_tweak(__m128i& cx, __m128& conc_var) _mm_store_si128((__m128i *)&l0[idx1 ^ 0x10], _mm_add_epi64(chunk3, bx1)); \ _mm_store_si128((__m128i *)&l0[idx1 ^ 0x20], _mm_add_epi64(chunk1, bx0)); \ _mm_store_si128((__m128i *)&l0[idx1 ^ 0x30], _mm_add_epi64(chunk2, ax0)); \ + if (ALGO == cryptonight_r) \ + cx = _mm_xor_si128(_mm_xor_si128(cx, chunk3), _mm_xor_si128(chunk1, chunk2)); \ } #define CN_MONERO_V8_SHUFFLE_1(n, l0, idx0, ax0, bx0, bx1, lo, hi) \ /* Shuffle the other 3x16 byte chunks in the current 64-byte cache line */ \ - if(ALGO == cryptonight_monero_v8) \ + if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_r_wow) \ { \ const uint64_t idx1 = idx0 & MASK; \ const __m128i chunk1 = _mm_xor_si128(_mm_load_si128((__m128i *)&l0[idx1 ^ 0x10]), _mm_set_epi64x(lo, hi)); \ @@ -667,6 +670,23 @@ inline void cryptonight_conceal_tweak(__m128i& cx, __m128& conc_var) assign(sqrt_result, int_sqrt33_1_double_precision(cx_64 + division_result)); \ } +#define CN_R_RANDOM_MATH(n, al, ah, cl, bx0, bx1, cn_r_data) \ + if (ALGO == cryptonight_r || ALGO == cryptonight_r_wow) \ + { \ + cl ^= (cn_r_data[0] + cn_r_data[1]) | ((uint64_t)(cn_r_data[2] + cn_r_data[3]) << 32); \ + cn_r_data[4] = static_cast(al); \ + cn_r_data[5] = static_cast(ah); \ + cn_r_data[6] = static_cast(_mm_cvtsi128_si32(bx0)); \ + cn_r_data[7] = static_cast(_mm_cvtsi128_si32(bx1)); \ + cn_r_data[8] = static_cast(_mm_cvtsi128_si32(_mm_srli_si128(bx1, 8))); \ + v4_random_math(ctx[n]->cn_r_ctx.code, cn_r_data); \ + } \ + if (ALGO == cryptonight_r) \ + { \ + al ^= cn_r_data[2] | ((uint64_t)(cn_r_data[3]) << 32); \ + ah ^= cn_r_data[0] | ((uint64_t)(cn_r_data[1]) << 32); \ + } + #define CN_INIT_SINGLE \ if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari || ALGO == cryptonight_bittube2) && len < 43) \ { \ @@ -674,7 +694,7 @@ inline void cryptonight_conceal_tweak(__m128i& cx, __m128& conc_var) return; \ } -#define CN_INIT(n, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, bx1, sqrt_result, division_result_xmm) \ +#define CN_INIT(n, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, bx1, sqrt_result, division_result_xmm, cn_r_data) \ keccak((const uint8_t *)input + len * n, len, ctx[n]->hash_state, 200); \ uint64_t monero_const; \ if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari || ALGO == cryptonight_bittube2) \ @@ -693,12 +713,13 @@ inline void cryptonight_conceal_tweak(__m128i& cx, __m128& conc_var) __m128i bx1; \ __m128i division_result_xmm; \ __m128 conc_var; \ - if(ALGO == cryptonight_conceal) \ + if(ALGO == cryptonight_conceal || ALGO == cryptonight_gpu) \ {\ - set_float_rounding_mode_conceal(); \ + set_float_rounding_mode_nearest(); \ conc_var = _mm_setzero_ps(); \ }\ GetOptimalSqrtType_t sqrt_result; \ + uint32_t cn_r_data[9]; \ /* END cryptonight_monero_v8 variables */ \ { \ uint64_t* h0 = (uint64_t*)ctx[n]->hash_state; \ @@ -712,6 +733,14 @@ inline void cryptonight_conceal_tweak(__m128i& cx, __m128& conc_var) assign(sqrt_result, h0[13]); \ set_float_rounding_mode(); \ } \ + if (ALGO == cryptonight_r || ALGO == cryptonight_r_wow) \ + { \ + bx1 = _mm_set_epi64x(h0[9] ^ h0[11], h0[8] ^ h0[10]); \ + cn_r_data[0] = (uint32_t)(h0[12]); \ + cn_r_data[1] = (uint32_t)(h0[12] >> 32); \ + cn_r_data[2] = (uint32_t)(h0[13]); \ + cn_r_data[3] = (uint32_t)(h0[13] >> 32); \ + } \ } \ __m128i *ptr0 @@ -732,7 +761,7 @@ inline void cryptonight_conceal_tweak(__m128i& cx, __m128& conc_var) else \ cx = _mm_aesenc_si128(cx, ax0); \ } \ - CN_MONERO_V8_SHUFFLE_0(n, l0, idx0, ax0, bx0, bx1) + CN_MONERO_V8_SHUFFLE_0(n, l0, idx0, ax0, bx0, bx1, cx) #define CN_STEP2(n, monero_const, l0, ax0, bx0, idx0, ptr0, cx) \ if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari || ALGO == cryptonight_bittube2) \ @@ -744,24 +773,32 @@ inline void cryptonight_conceal_tweak(__m128i& cx, __m128& conc_var) ptr0 = (__m128i *)&l0[idx0 & MASK]; \ if(PREFETCH) \ _mm_prefetch((const char*)ptr0, _MM_HINT_T0); \ - if(ALGO != cryptonight_monero_v8) \ + if(ALGO != cryptonight_monero_v8 && ALGO != cryptonight_r && ALGO != cryptonight_r_wow) \ bx0 = cx -#define CN_STEP3(n, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0, cx, bx1, sqrt_result, division_result_xmm) \ +#define CN_STEP3(n, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0, cx, bx1, sqrt_result, division_result_xmm, cn_r_data) \ uint64_t lo, cl, ch; \ uint64_t al0 = _mm_cvtsi128_si64(ax0); \ uint64_t ah0 = ((uint64_t*)&ax0)[1]; \ cl = ((uint64_t*)ptr0)[0]; \ ch = ((uint64_t*)ptr0)[1]; \ + CN_R_RANDOM_MATH(n, al0, ah0, cl, bx0, bx1, cn_r_data); \ CN_MONERO_V8_DIV(n, cx, sqrt_result, division_result_xmm, cl); \ { \ uint64_t hi; \ lo = _umul128(idx0, cl, &hi); \ - CN_MONERO_V8_SHUFFLE_1(n, l0, idx0, ax0, bx0, bx1, lo, hi); \ + if(ALGO == cryptonight_r) \ + { \ + CN_MONERO_V8_SHUFFLE_0(n, l0, idx0, ax0, bx0, bx1, cx); \ + } \ + else \ + { \ + CN_MONERO_V8_SHUFFLE_1(n, l0, idx0, ax0, bx0, bx1, lo, hi); \ + } \ ah0 += lo; \ al0 += hi; \ } \ - if(ALGO == cryptonight_monero_v8) \ + if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_r || ALGO != cryptonight_r_wow) \ { \ bx1 = bx0; \ bx0 = cx; \ @@ -851,6 +888,7 @@ inline void cryptonight_conceal_tweak(__m128i& cx, __m128& conc_var) #define CN_ENUM_13(n, x1, x2, x3, x4, x5, x6, x7, x8, x9, x10, x11, x12, x13) n, x1 ## n, x2 ## n, x3 ## n, x4 ## n, x5 ## n, x6 ## n, x7 ## n, x8 ## n, x9 ## n, x10 ## n, x11 ## n, x12 ## n, x13 ## n #define CN_ENUM_14(n, x1, x2, x3, x4, x5, x6, x7, x8, x9, x10, x11, x12, x13, x14) n, x1 ## n, x2 ## n, x3 ## n, x4 ## n, x5 ## n, x6 ## n, x7 ## n, x8 ## n, x9 ## n, x10 ## n, x11 ## n, x12 ## n, x13 ## n, x14 ## n #define CN_ENUM_15(n, x1, x2, x3, x4, x5, x6, x7, x8, x9, x10, x11, x12, x13, x14, x15) n, x1 ## n, x2 ## n, x3 ## n, x4 ## n, x5 ## n, x6 ## n, x7 ## n, x8 ## n, x9 ## n, x10 ## n, x11 ## n, x12 ## n, x13 ## n, x14 ## n, x15 ## n +#define CN_ENUM_16(n, x1, x2, x3, x4, x5, x6, x7, x8, x9, x10, x11, x12, x13, x14, x15, x16) n, x1 ## n, x2 ## n, x3 ## n, x4 ## n, x5 ## n, x6 ## n, x7 ## n, x8 ## n, x9 ## n, x10 ## n, x11 ## n, x12 ## n, x13 ## n, x14 ## n, x15 ## n, x16 ## n /** repeat a macro call multiple times * @@ -886,14 +924,14 @@ struct Cryptonight_hash<1> const size_t MEM = algo.Mem(); CN_INIT_SINGLE; - REPEAT_1(10, CN_INIT, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, bx1, sqrt_result, division_result_xmm); + REPEAT_1(11, CN_INIT, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, bx1, sqrt_result, division_result_xmm, cn_r_data); // Optim - 90% time boundary for(size_t i = 0; i < ITERATIONS; i++) { REPEAT_1(9, CN_STEP1, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, cx, bx1); REPEAT_1(7, CN_STEP2, monero_const, l0, ax0, bx0, idx0, ptr0, cx); - REPEAT_1(15, CN_STEP3, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0, cx, bx1, sqrt_result, division_result_xmm); + REPEAT_1(16, CN_STEP3, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0, cx, bx1, sqrt_result, division_result_xmm, cn_r_data); REPEAT_1(11, CN_STEP4, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0); REPEAT_1(6, CN_STEP5, monero_const, l0, ax0, bx0, idx0, ptr0); } @@ -915,14 +953,14 @@ struct Cryptonight_hash<2> const size_t MEM = algo.Mem(); CN_INIT_SINGLE; - REPEAT_2(10, CN_INIT, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, bx1, sqrt_result, division_result_xmm); + REPEAT_2(11, CN_INIT, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, bx1, sqrt_result, division_result_xmm, cn_r_data); // Optim - 90% time boundary for(size_t i = 0; i < ITERATIONS; i++) { REPEAT_2(9, CN_STEP1, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, cx, bx1); REPEAT_2(7, CN_STEP2, monero_const, l0, ax0, bx0, idx0, ptr0, cx); - REPEAT_2(15, CN_STEP3, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0, cx, bx1, sqrt_result, division_result_xmm); + REPEAT_2(16, CN_STEP3, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0, cx, bx1, sqrt_result, division_result_xmm, cn_r_data); REPEAT_2(11, CN_STEP4, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0); REPEAT_2(6, CN_STEP5, monero_const, l0, ax0, bx0, idx0, ptr0); } @@ -944,14 +982,14 @@ struct Cryptonight_hash<3> const size_t MEM = algo.Mem(); CN_INIT_SINGLE; - REPEAT_3(10, CN_INIT, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, bx1, sqrt_result, division_result_xmm); + REPEAT_3(11, CN_INIT, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, bx1, sqrt_result, division_result_xmm, cn_r_data); // Optim - 90% time boundary for(size_t i = 0; i < ITERATIONS; i++) { REPEAT_3(9, CN_STEP1, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, cx, bx1); REPEAT_3(7, CN_STEP2, monero_const, l0, ax0, bx0, idx0, ptr0, cx); - REPEAT_3(15, CN_STEP3, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0, cx, bx1, sqrt_result, division_result_xmm); + REPEAT_3(16, CN_STEP3, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0, cx, bx1, sqrt_result, division_result_xmm, cn_r_data); REPEAT_3(11, CN_STEP4, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0); REPEAT_3(6, CN_STEP5, monero_const, l0, ax0, bx0, idx0, ptr0); } @@ -973,14 +1011,14 @@ struct Cryptonight_hash<4> const size_t MEM = algo.Mem(); CN_INIT_SINGLE; - REPEAT_4(10, CN_INIT, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, bx1, sqrt_result, division_result_xmm); + REPEAT_4(11, CN_INIT, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, bx1, sqrt_result, division_result_xmm, cn_r_data); // Optim - 90% time boundary for(size_t i = 0; i < ITERATIONS; i++) { REPEAT_4(9, CN_STEP1, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, cx, bx1); REPEAT_4(7, CN_STEP2, monero_const, l0, ax0, bx0, idx0, ptr0, cx); - REPEAT_4(15, CN_STEP3, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0, cx, bx1, sqrt_result, division_result_xmm); + REPEAT_4(16, CN_STEP3, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0, cx, bx1, sqrt_result, division_result_xmm, cn_r_data); REPEAT_4(11, CN_STEP4, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0); REPEAT_4(6, CN_STEP5, monero_const, l0, ax0, bx0, idx0, ptr0); } @@ -1002,14 +1040,14 @@ struct Cryptonight_hash<5> const size_t MEM = algo.Mem(); CN_INIT_SINGLE; - REPEAT_5(10, CN_INIT, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, bx1, sqrt_result, division_result_xmm); + REPEAT_5(11, CN_INIT, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, bx1, sqrt_result, division_result_xmm, cn_r_data); // Optim - 90% time boundary for(size_t i = 0; i < ITERATIONS; i++) { REPEAT_5(9, CN_STEP1, monero_const, conc_var, l0, ax0, bx0, idx0, ptr0, cx, bx1); REPEAT_5(7, CN_STEP2, monero_const, l0, ax0, bx0, idx0, ptr0, cx); - REPEAT_5(15, CN_STEP3, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0, cx, bx1, sqrt_result, division_result_xmm); + REPEAT_5(16, CN_STEP3, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0, cx, bx1, sqrt_result, division_result_xmm, cn_r_data); REPEAT_5(11, CN_STEP4, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0); REPEAT_5(6, CN_STEP5, monero_const, l0, ax0, bx0, idx0, ptr0); } @@ -1034,7 +1072,6 @@ struct Cryptonight_hash_asm<1, asm_version> template static void hash(const void* input, size_t len, void* output, cryptonight_ctx** ctx, const xmrstak_algo& algo) { - keccak((const uint8_t *)input, len, ctx[0]->hash_state, 200); cn_explode_scratchpad((__m128i*)ctx[0]->hash_state, (__m128i*)ctx[0]->long_state, algo); @@ -1087,7 +1124,6 @@ struct Cryptonight_hash_gpu template static void hash(const void* input, size_t len, void* output, cryptonight_ctx** ctx, const xmrstak_algo& algo) { - keccak((const uint8_t *)input, len, ctx[0]->hash_state, 200); cn_explode_scratchpad_gpu(ctx[0]->hash_state, ctx[0]->long_state, algo); @@ -1101,3 +1137,20 @@ struct Cryptonight_hash_gpu memcpy(output, ctx[0]->hash_state, 32); } }; + +template +struct Cryptonight_R_generator +{ + template + static void cn_on_new_job(const xmrstak::miner_work& work, cryptonight_ctx** ctx) + { + if(ctx[0]->cn_r_ctx.height == work.iBlockHeight) + return; + + ctx[0]->cn_r_ctx.height = work.iBlockHeight; + v4_random_math_init(ctx[0]->cn_r_ctx.code, work.iBlockHeight); + + for(size_t i=1; i < N; i++) + ctx[i]->cn_r_ctx = ctx[0]->cn_r_ctx; + } +}; diff --git a/xmrstak/backend/cpu/crypto/variant4_random_math.h b/xmrstak/backend/cpu/crypto/variant4_random_math.h new file mode 100644 index 000000000..07dd3cf61 --- /dev/null +++ b/xmrstak/backend/cpu/crypto/variant4_random_math.h @@ -0,0 +1,451 @@ +#pragma once + +#include +#include "../../cryptonight.hpp" + +extern "C" +{ + #include "c_blake256.h" +} + +enum V4_Settings +{ + // Generate code with minimal theoretical latency = 45 cycles, which is equivalent to 15 multiplications + TOTAL_LATENCY = 15 * 3, + + // Always generate at least 60 instructions + NUM_INSTRUCTIONS_MIN = 60, + + // Never generate more than 70 instructions (final RET instruction doesn't count here) + NUM_INSTRUCTIONS_MAX = 70, + + // Available ALUs for MUL + // Modern CPUs typically have only 1 ALU which can do multiplications + ALU_COUNT_MUL = 1, + + // Total available ALUs + // Modern CPUs have 4 ALUs, but we use only 3 because random math executes together with other main loop code + ALU_COUNT = 3, +}; + +enum V4_InstructionList +{ + MUL, // a*b + ADD, // a+b + C, C is an unsigned 32-bit constant + SUB, // a-b + ROR, // rotate right "a" by "b & 31" bits + ROL, // rotate left "a" by "b & 31" bits + XOR, // a^b + RET, // finish execution + V4_INSTRUCTION_COUNT = RET, +}; + +// V4_InstructionDefinition is used to generate code from random data +// Every random sequence of bytes is a valid code +// +// There are 9 registers in total: +// - 4 variable registers +// - 5 constant registers initialized from loop variables +// This is why dst_index is 2 bits +enum V4_InstructionDefinition +{ + V4_OPCODE_BITS = 3, + V4_DST_INDEX_BITS = 2, + V4_SRC_INDEX_BITS = 3, +}; + +struct V4_Instruction +{ + uint8_t opcode; + uint8_t dst_index; + uint8_t src_index; + uint32_t C; +}; + +#ifndef FORCEINLINE +#ifdef __GNUC__ +#define FORCEINLINE __attribute__((always_inline)) inline +#elif _MSC_VER +#define FORCEINLINE __forceinline +#else +#define FORCEINLINE inline +#endif +#endif + +#ifndef UNREACHABLE_CODE +#ifdef __GNUC__ +#define UNREACHABLE_CODE __builtin_unreachable() +#elif _MSC_VER +#define UNREACHABLE_CODE __assume(false) +#else +#define UNREACHABLE_CODE +#endif +#endif + +// Random math interpreter's loop is fully unrolled and inlined to achieve 100% branch prediction on CPU: +// every switch-case will point to the same destination on every iteration of Cryptonight main loop +// +// This is about as fast as it can get without using low-level machine code generation +template +static void v4_random_math(const struct V4_Instruction* code, v4_reg* r) +{ + enum + { + REG_BITS = sizeof(v4_reg) * 8, + }; + +#define V4_EXEC(i) \ + { \ + const struct V4_Instruction* op = code + i; \ + const v4_reg src = r[op->src_index]; \ + v4_reg* dst = r + op->dst_index; \ + switch (op->opcode) \ + { \ + case MUL: \ + *dst *= src; \ + break; \ + case ADD: \ + *dst += src + op->C; \ + break; \ + case SUB: \ + *dst -= src; \ + break; \ + case ROR: \ + { \ + const uint32_t shift = src % REG_BITS; \ + *dst = (*dst >> shift) | (*dst << ((REG_BITS - shift) % REG_BITS)); \ + } \ + break; \ + case ROL: \ + { \ + const uint32_t shift = src % REG_BITS; \ + *dst = (*dst << shift) | (*dst >> ((REG_BITS - shift) % REG_BITS)); \ + } \ + break; \ + case XOR: \ + *dst ^= src; \ + break; \ + case RET: \ + return; \ + default: \ + UNREACHABLE_CODE; \ + break; \ + } \ + } + +#define V4_EXEC_10(j) \ + V4_EXEC(j + 0) \ + V4_EXEC(j + 1) \ + V4_EXEC(j + 2) \ + V4_EXEC(j + 3) \ + V4_EXEC(j + 4) \ + V4_EXEC(j + 5) \ + V4_EXEC(j + 6) \ + V4_EXEC(j + 7) \ + V4_EXEC(j + 8) \ + V4_EXEC(j + 9) + + // Generated program can have 60 + a few more (usually 2-3) instructions to achieve required latency + // I've checked all block heights < 10,000,000 and here is the distribution of program sizes: + // + // 60 27960 + // 61 105054 + // 62 2452759 + // 63 5115997 + // 64 1022269 + // 65 1109635 + // 66 153145 + // 67 8550 + // 68 4529 + // 69 102 + + // Unroll 70 instructions here + V4_EXEC_10(0); // instructions 0-9 + V4_EXEC_10(10); // instructions 10-19 + V4_EXEC_10(20); // instructions 20-29 + V4_EXEC_10(30); // instructions 30-39 + V4_EXEC_10(40); // instructions 40-49 + V4_EXEC_10(50); // instructions 50-59 + V4_EXEC_10(60); // instructions 60-69 + +#undef V4_EXEC_10 +#undef V4_EXEC +} + +// If we don't have enough data available, generate more +static FORCEINLINE void check_data(size_t* data_index, const size_t bytes_needed, int8_t* data, const size_t data_size) +{ + if (*data_index + bytes_needed > data_size) + { + blake256_hash((uint8_t*)data, (uint8_t*)data, data_size); + *data_index = 0; + } +} + +#define SWAP32LE(x) x +#define SWAP64LE(x) x + +// Generates as many random math operations as possible with given latency and ALU restrictions +// "code" array must have space for NUM_INSTRUCTIONS_MAX+1 instructions +template +static int v4_random_math_init(struct V4_Instruction* code, const uint64_t height) +{ + // MUL is 3 cycles, 3-way addition and rotations are 2 cycles, SUB/XOR are 1 cycle + // These latencies match real-life instruction latencies for Intel CPUs starting from Sandy Bridge and up to Skylake/Coffee lake + // + // AMD Ryzen has the same latencies except 1-cycle ROR/ROL, so it'll be a bit faster than Intel Sandy Bridge and newer processors + // Surprisingly, Intel Nehalem also has 1-cycle ROR/ROL, so it'll also be faster than Intel Sandy Bridge and newer processors + // AMD Bulldozer has 4 cycles latency for MUL (slower than Intel) and 1 cycle for ROR/ROL (faster than Intel), so average performance will be the same + // Source: https://www.agner.org/optimize/instruction_tables.pdf + const int op_latency[V4_INSTRUCTION_COUNT] = { 3, 2, 1, 2, 2, 1 }; + + // Instruction latencies for theoretical ASIC implementation + const int asic_op_latency[V4_INSTRUCTION_COUNT] = { 3, 1, 1, 1, 1, 1 }; + + // Available ALUs for each instruction + const int op_ALUs[V4_INSTRUCTION_COUNT] = { ALU_COUNT_MUL, ALU_COUNT, ALU_COUNT, ALU_COUNT, ALU_COUNT, ALU_COUNT }; + + int8_t data[32]; + memset(data, 0, sizeof(data)); + uint64_t tmp = SWAP64LE(height); + memcpy(data, &tmp, sizeof(uint64_t)); + if(ALGO == cryptonight_r) + { + data[20] = -38; + } + + // Set data_index past the last byte in data + // to trigger full data update with blake hash + // before we start using it + size_t data_index = sizeof(data); + + int code_size; + + // There is a small chance (1.8%) that register R8 won't be used in the generated program + // So we keep track of it and try again if it's not used + bool r8_used; + do { + int latency[9]; + int asic_latency[9]; + + // Tracks previous instruction and value of the source operand for registers R0-R3 throughout code execution + // byte 0: current value of the destination register + // byte 1: instruction opcode + // byte 2: current value of the source register + // + // Registers R4-R8 are constant and are treated as having the same value because when we do + // the same operation twice with two constant source registers, it can be optimized into a single operation + uint32_t inst_data[9] = { 0, 1, 2, 3, 0xFFFFFF, 0xFFFFFF, 0xFFFFFF, 0xFFFFFF, 0xFFFFFF }; + + bool alu_busy[TOTAL_LATENCY + 1][ALU_COUNT]; + bool is_rotation[V4_INSTRUCTION_COUNT]; + bool rotated[4]; + int rotate_count = 0; + + memset(latency, 0, sizeof(latency)); + memset(asic_latency, 0, sizeof(asic_latency)); + memset(alu_busy, 0, sizeof(alu_busy)); + memset(is_rotation, 0, sizeof(is_rotation)); + memset(rotated, 0, sizeof(rotated)); + is_rotation[ROR] = true; + is_rotation[ROL] = true; + + int num_retries = 0; + code_size = 0; + + int total_iterations = 0; + r8_used = (ALGO == cryptonight_r_wow); + + // Generate random code to achieve minimal required latency for our abstract CPU + // Try to get this latency for all 4 registers + while (((latency[0] < TOTAL_LATENCY) || (latency[1] < TOTAL_LATENCY) || (latency[2] < TOTAL_LATENCY) || (latency[3] < TOTAL_LATENCY)) && (num_retries < 64)) + { + // Fail-safe to guarantee loop termination + ++total_iterations; + if (total_iterations > 256) + break; + + check_data(&data_index, 1, data, sizeof(data)); + + const uint8_t c = ((uint8_t*)data)[data_index++]; + + // MUL = opcodes 0-2 + // ADD = opcode 3 + // SUB = opcode 4 + // ROR/ROL = opcode 5, shift direction is selected randomly + // XOR = opcodes 6-7 + uint8_t opcode = c & ((1 << V4_OPCODE_BITS) - 1); + if (opcode == 5) + { + check_data(&data_index, 1, data, sizeof(data)); + opcode = (data[data_index++] >= 0) ? ROR : ROL; + } + else if (opcode >= 6) + { + opcode = XOR; + } + else + { + opcode = (opcode <= 2) ? MUL : (opcode - 2); + } + + uint8_t dst_index = (c >> V4_OPCODE_BITS) & ((1 << V4_DST_INDEX_BITS) - 1); + uint8_t src_index = (c >> (V4_OPCODE_BITS + V4_DST_INDEX_BITS)) & ((1 << V4_SRC_INDEX_BITS) - 1); + + const int a = dst_index; + int b = src_index; + + // Don't do ADD/SUB/XOR with the same register + if (((opcode == ADD) || (opcode == SUB) || (opcode == XOR)) && (a == b)) + { + // a is always < 4, so we don't need to check bounds here + b = (ALGO == cryptonight_r_wow) ? (a + 4) : 8; + src_index = b; + } + + // Don't do rotation with the same destination twice because it's equal to a single rotation + if (is_rotation[opcode] && rotated[a]) + { + continue; + } + + // Don't do the same instruction (except MUL) with the same source value twice because all other cases can be optimized: + // 2xADD(a, b, C) = ADD(a, b*2, C1+C2), same for SUB and rotations + // 2xXOR(a, b) = NOP + if ((opcode != MUL) && ((inst_data[a] & 0xFFFF00) == (opcode << 8) + ((inst_data[b] & 255) << 16))) + { + continue; + } + + // Find which ALU is available (and when) for this instruction + int next_latency = (latency[a] > latency[b]) ? latency[a] : latency[b]; + int alu_index = -1; + while (next_latency < TOTAL_LATENCY) + { + for (int i = op_ALUs[opcode] - 1; i >= 0; --i) + { + if (!alu_busy[next_latency][i]) + { + // ADD is implemented as two 1-cycle instructions on a real CPU, so do an additional availability check + if ((opcode == ADD) && alu_busy[next_latency + 1][i]) + { + continue; + } + + // Rotation can only start when previous rotation is finished, so do an additional availability check + if (is_rotation[opcode] && (next_latency < rotate_count * op_latency[opcode])) + { + continue; + } + + alu_index = i; + break; + } + } + if (alu_index >= 0) + { + break; + } + ++next_latency; + } + + // Don't generate instructions that leave some register unchanged for more than 7 cycles + if (next_latency > latency[a] + 7) + { + continue; + } + + next_latency += op_latency[opcode]; + + if (next_latency <= TOTAL_LATENCY) + { + if (is_rotation[opcode]) + { + ++rotate_count; + } + + // Mark ALU as busy only for the first cycle when it starts executing the instruction because ALUs are fully pipelined + alu_busy[next_latency - op_latency[opcode]][alu_index] = true; + latency[a] = next_latency; + + // ASIC is supposed to have enough ALUs to run as many independent instructions per cycle as possible, so latency calculation for ASIC is simple + asic_latency[a] = ((asic_latency[a] > asic_latency[b]) ? asic_latency[a] : asic_latency[b]) + asic_op_latency[opcode]; + + rotated[a] = is_rotation[opcode]; + + inst_data[a] = code_size + (opcode << 8) + ((inst_data[b] & 255) << 16); + + code[code_size].opcode = opcode; + code[code_size].dst_index = dst_index; + code[code_size].src_index = src_index; + code[code_size].C = 0; + + if (src_index == 8) + { + r8_used = true; + } + + if (opcode == ADD) + { + // ADD instruction is implemented as two 1-cycle instructions on a real CPU, so mark ALU as busy for the next cycle too + alu_busy[next_latency - op_latency[opcode] + 1][alu_index] = true; + + // ADD instruction requires 4 more random bytes for 32-bit constant "C" in "a = a + b + C" + check_data(&data_index, sizeof(uint32_t), data, sizeof(data)); + uint32_t t; + memcpy(&t, data + data_index, sizeof(uint32_t)); + code[code_size].C = SWAP32LE(t); + data_index += sizeof(uint32_t); + } + + ++code_size; + if (code_size >= NUM_INSTRUCTIONS_MIN) + { + break; + } + } + else + { + ++num_retries; + } + } + + // ASIC has more execution resources and can extract as much parallelism from the code as possible + // We need to add a few more MUL and ROR instructions to achieve minimal required latency for ASIC + // Get this latency for at least 1 of the 4 registers + const int prev_code_size = code_size; + while ((code_size < NUM_INSTRUCTIONS_MAX) && (asic_latency[0] < TOTAL_LATENCY) && (asic_latency[1] < TOTAL_LATENCY) && (asic_latency[2] < TOTAL_LATENCY) && (asic_latency[3] < TOTAL_LATENCY)) + { + int min_idx = 0; + int max_idx = 0; + for (int i = 1; i < 4; ++i) + { + if (asic_latency[i] < asic_latency[min_idx]) min_idx = i; + if (asic_latency[i] > asic_latency[max_idx]) max_idx = i; + } + + const uint8_t pattern[3] = { ROR, MUL, MUL }; + const uint8_t opcode = pattern[(code_size - prev_code_size) % 3]; + latency[min_idx] = latency[max_idx] + op_latency[opcode]; + asic_latency[min_idx] = asic_latency[max_idx] + asic_op_latency[opcode]; + + code[code_size].opcode = opcode; + code[code_size].dst_index = min_idx; + code[code_size].src_index = max_idx; + code[code_size].C = 0; + ++code_size; + } + + // There is ~98.15% chance that loop condition is false, so this loop will execute only 1 iteration most of the time + // It never does more than 4 iterations for all block heights < 10,000,000 + } while (!r8_used || (code_size < NUM_INSTRUCTIONS_MIN) || (code_size > NUM_INSTRUCTIONS_MAX)); + + // It's guaranteed that NUM_INSTRUCTIONS_MIN <= code_size <= NUM_INSTRUCTIONS_MAX here + // Add final instruction to stop the interpreter + code[code_size].opcode = RET; + code[code_size].dst_index = 0; + code[code_size].src_index = 0; + code[code_size].C = 0; + + return code_size; +} diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp index 50507f2ae..064b07339 100644 --- a/xmrstak/backend/cpu/minethd.cpp +++ b/xmrstak/backend/cpu/minethd.cpp @@ -50,6 +50,7 @@ #include #include #include +#include #ifdef _WIN32 #include @@ -250,34 +251,35 @@ bool minethd::self_test() hashf("This is a test", 14, out, ctx, algo); bResult = bResult && memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 32) == 0; + minethd::cn_on_new_job dm; hashf = func_selector(::jconf::inst()->HaveHardwareAes(), true, algo); hashf("This is a test", 14, out, ctx, algo); bResult = bResult && memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 32) == 0; - hashf_multi = func_multi_selector<2>(::jconf::inst()->HaveHardwareAes(), false, algo); + func_multi_selector<2>(hashf_multi, dm, ::jconf::inst()->HaveHardwareAes(), false, algo); hashf_multi("The quick brown fox jumps over the lazy dogThe quick brown fox jumps over the lazy log", 43, out, ctx, algo); bResult = bResult && memcmp(out, "\x3e\xbb\x7f\x9f\x7d\x27\x3d\x7c\x31\x8d\x86\x94\x77\x55\x0c\xc8\x00\xcf\xb1\x1b\x0c\xad\xb7\xff\xbd\xf6\xf8\x9f\x3a\x47\x1c\x59" "\xb4\x77\xd5\x02\xe4\xd8\x48\x7f\x42\xdf\xe3\x8e\xed\x73\x81\x7a\xda\x91\xb7\xe2\x63\xd2\x91\x71\xb6\x5c\x44\x3a\x01\x2a\x41\x22", 64) == 0; - hashf_multi = func_multi_selector<2>(::jconf::inst()->HaveHardwareAes(), true, algo); + func_multi_selector<2>(hashf_multi, dm, ::jconf::inst()->HaveHardwareAes(), true, algo); hashf_multi("The quick brown fox jumps over the lazy dogThe quick brown fox jumps over the lazy log", 43, out, ctx, algo); bResult = bResult && memcmp(out, "\x3e\xbb\x7f\x9f\x7d\x27\x3d\x7c\x31\x8d\x86\x94\x77\x55\x0c\xc8\x00\xcf\xb1\x1b\x0c\xad\xb7\xff\xbd\xf6\xf8\x9f\x3a\x47\x1c\x59" "\xb4\x77\xd5\x02\xe4\xd8\x48\x7f\x42\xdf\xe3\x8e\xed\x73\x81\x7a\xda\x91\xb7\xe2\x63\xd2\x91\x71\xb6\x5c\x44\x3a\x01\x2a\x41\x22", 64) == 0; - hashf_multi = func_multi_selector<3>(::jconf::inst()->HaveHardwareAes(), false, algo); + func_multi_selector<3>(hashf_multi, dm, ::jconf::inst()->HaveHardwareAes(), false, algo); hashf_multi("This is a testThis is a testThis is a test", 14, out, ctx, algo); bResult = bResult && memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05" "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05" "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 96) == 0; - hashf_multi = func_multi_selector<4>(::jconf::inst()->HaveHardwareAes(), false, algo); + func_multi_selector<4>(hashf_multi, dm, ::jconf::inst()->HaveHardwareAes(), false, algo); hashf_multi("This is a testThis is a testThis is a testThis is a test", 14, out, ctx, algo); bResult = bResult && memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05" "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05" "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05" "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 128) == 0; - hashf_multi = func_multi_selector<5>(::jconf::inst()->HaveHardwareAes(), false, algo); + func_multi_selector<5>(hashf_multi, dm, ::jconf::inst()->HaveHardwareAes(), false, algo); hashf_multi("This is a testThis is a testThis is a testThis is a testThis is a test", 14, out, ctx, algo); bResult = bResult && memcmp(out, "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05" "\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05" @@ -427,6 +429,16 @@ bool minethd::self_test() hashf("This is a test This is a test This is a test", 44, out, ctx, algo); bResult = bResult && memcmp(out, "\x30\x5f\x66\xfe\xbb\xf3\x60\x0e\xda\xbb\x60\xf7\xf1\xc9\xb9\x0a\x3a\xe8\x5a\x31\xd4\x76\xca\x38\x1d\x56\x18\xa6\xc6\x27\x60\xd7", 32) == 0; } + else if(algo == POW(cryptonight_r)) + { + minethd::cn_on_new_job set_job; + func_multi_selector<1>(hashf, set_job, ::jconf::inst()->HaveHardwareAes(), false, algo); + miner_work work; + work.iBlockHeight = 1806260; + set_job(work, ctx); + hashf("\x54\x68\x69\x73\x20\x69\x73\x20\x61\x20\x74\x65\x73\x74\x20\x54\x68\x69\x73\x20\x69\x73\x20\x61\x20\x74\x65\x73\x74\x20\x54\x68\x69\x73\x20\x69\x73\x20\x61\x20\x74\x65\x73\x74", 44, out, ctx, algo); + bResult = bResult && memcmp(out, "\xf7\x59\x58\x8a\xd5\x7e\x75\x84\x67\x29\x54\x43\xa9\xbd\x71\x49\x0a\xbf\xf8\xe9\xda\xd1\xb9\x5b\x6b\xf2\xf5\xd0\xd7\x83\x87\xbc", 32) == 0; + } else printer::inst()->print_msg(L0, "Cryptonight hash self-test NOT defined for POW %s", algo.Name().c_str()); @@ -511,7 +523,8 @@ static std::string getAsmName(const uint32_t num_hashes) } template -minethd::cn_hash_fun minethd::func_multi_selector(bool bHaveAes, bool bNoPrefetch, const xmrstak_algo& algo, const std::string& asm_version_str) +void minethd::func_multi_selector(minethd::cn_hash_fun& hash_fun, minethd::cn_on_new_job& on_new_job, + bool bHaveAes, bool bNoPrefetch, const xmrstak_algo& algo, const std::string& asm_version_str) { static_assert(N >= 1, "number of threads must be >= 1" ); @@ -564,6 +577,9 @@ minethd::cn_hash_fun minethd::func_multi_selector(bool bHaveAes, bool bNoPrefetc case cryptonight_conceal: algv = 13; break; + case cryptonight_r: + algv = 14; + break; default: algv = 2; break; @@ -638,15 +654,19 @@ minethd::cn_hash_fun minethd::func_multi_selector(bool bHaveAes, bool bNoPrefetc Cryptonight_hash::template hash, Cryptonight_hash::template hash, Cryptonight_hash::template hash, - Cryptonight_hash::template hash + Cryptonight_hash::template hash, + + Cryptonight_hash::template hash, + Cryptonight_hash::template hash, + Cryptonight_hash::template hash, + Cryptonight_hash::template hash }; std::bitset<2> digit; digit.set(0, !bHaveAes); digit.set(1, !bNoPrefetch); - auto selected_function = func_table[ algv << 2 | digit.to_ulong() ]; - + hash_fun = func_table[ algv << 2 | digit.to_ulong() ]; // check for asm optimized version for cryptonight_v8 if(N <= 2 && algo == cryptonight_monero_v8 && bHaveAes && algo.Mem() == CN_MEMORY && algo.Iter() == CN_ITER) @@ -661,15 +681,15 @@ minethd::cn_hash_fun minethd::func_multi_selector(bool bHaveAes, bool bNoPrefetc { // Intel Ivy Bridge (Xeon v2, Core i7/i5/i3 3xxx, Pentium G2xxx, Celeron G1xxx) if(N == 1) - selected_function = Cryptonight_hash_asm<1u, 0u>::template hash; + hash_fun = Cryptonight_hash_asm<1u, 0u>::template hash; else if(N == 2) - selected_function = Cryptonight_hash_asm<2u, 0u>::template hash; + hash_fun = Cryptonight_hash_asm<2u, 0u>::template hash; } // supports only 1 thread per hash if(N == 1 && selected_asm == "amd_avx") { // AMD Ryzen (1xxx and 2xxx series) - selected_function = Cryptonight_hash_asm<1u, 1u>::template hash; + hash_fun = Cryptonight_hash_asm<1u, 1u>::template hash; } if(asm_version_str == "auto" && (selected_asm != "intel_avx" || selected_asm != "amd_avx")) printer::inst()->print_msg(L3, "Switch to assembler version for '%s' cpu's", selected_asm.c_str()); @@ -678,12 +698,23 @@ minethd::cn_hash_fun minethd::func_multi_selector(bool bHaveAes, bool bNoPrefetc } } - return selected_function; + static const std::unordered_map on_new_job_map = { + {cryptonight_r, Cryptonight_R_generator::template cn_on_new_job}, + }; + + auto it = on_new_job_map.find(algo.Id()); + if (it != on_new_job_map.end()) + on_new_job = it->second; + else + on_new_job = nullptr; } minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, const xmrstak_algo& algo) { - return func_multi_selector<1>(bHaveAes, bNoPrefetch, algo); + minethd::cn_hash_fun fun; + minethd::cn_on_new_job dm; + func_multi_selector<1>(fun, dm, bHaveAes, bNoPrefetch, algo); + return fun; } void minethd::work_main() @@ -763,10 +794,12 @@ void minethd::multiway_work_main() // start with root algorithm and switch later if fork version is reached auto miner_algo = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot(); - cn_hash_fun hash_fun_multi = func_multi_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo, asm_version_str); + cn_hash_fun hash_fun_multi; + cn_on_new_job on_new_job; uint8_t version = 0; size_t lastPoolId = 0; + func_multi_selector(hash_fun_multi, on_new_job, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo, asm_version_str); while (bQuit == 0) { if (oWork.bStall) @@ -798,17 +831,20 @@ void minethd::multiway_work_main() if(new_version >= coinDesc.GetMiningForkVersion()) { miner_algo = coinDesc.GetMiningAlgo(); - hash_fun_multi = func_multi_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo, asm_version_str); + func_multi_selector(hash_fun_multi, on_new_job, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo, asm_version_str); } else { miner_algo = coinDesc.GetMiningAlgoRoot(); - hash_fun_multi = func_multi_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo, asm_version_str); + func_multi_selector(hash_fun_multi, on_new_job, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo, asm_version_str); } lastPoolId = oWork.iPoolId; version = new_version; } + if(on_new_job != nullptr) + on_new_job(oWork, ctx); + while (globalStates::inst().iGlobalJobNo.load(std::memory_order_relaxed) == iJobNo) { if ((iCount++ & 0x7) == 0) //Store stats every 8*N hashes diff --git a/xmrstak/backend/cpu/minethd.hpp b/xmrstak/backend/cpu/minethd.hpp index 41315269b..ca89e5b52 100644 --- a/xmrstak/backend/cpu/minethd.hpp +++ b/xmrstak/backend/cpu/minethd.hpp @@ -22,6 +22,7 @@ class minethd : public iBackend static std::vector thread_starter(uint32_t threadOffset, miner_work& pWork); static bool self_test(); + typedef void (*cn_on_new_job)(const miner_work&, cryptonight_ctx**); typedef void (*cn_hash_fun)(const void*, size_t, void*, cryptonight_ctx**, const xmrstak_algo&); static cn_hash_fun func_selector(bool bHaveAes, bool bNoPrefetch, const xmrstak_algo& algo); @@ -29,11 +30,12 @@ class minethd : public iBackend static cryptonight_ctx* minethd_alloc_ctx(); -private: - template - static cn_hash_fun func_multi_selector(bool bHaveAes, bool bNoPrefetch, const xmrstak_algo& algo, const std::string& asm_version_str = "off"); + static void func_multi_selector(minethd::cn_hash_fun& hash_fun, minethd::cn_on_new_job& on_new_job, + bool bHaveAes, bool bNoPrefetch, const xmrstak_algo& algo, const std::string& asm_version_str = "off"); + private: + minethd(miner_work& pWork, size_t iNo, int iMultiway, bool no_prefetch, int64_t affinity, const std::string& asm_version); template diff --git a/xmrstak/backend/cryptonight.hpp b/xmrstak/backend/cryptonight.hpp index db0787789..00311bb93 100644 --- a/xmrstak/backend/cryptonight.hpp +++ b/xmrstak/backend/cryptonight.hpp @@ -24,6 +24,8 @@ enum xmrstak_algo_id cryptonight_superfast = 12, cryptonight_gpu = 13, cryptonight_conceal = 14, + cryptonight_r_wow = 15, + cryptonight_r = 16, cryptonight_turtle = start_derived_algo_id, cryptonight_v8_half = (start_derived_algo_id + 1), @@ -37,7 +39,7 @@ enum xmrstak_algo_id */ inline std::string get_algo_name(xmrstak_algo_id algo_id) { - static std::array base_algo_names = + static std::array base_algo_names = {{ "invalid_algo", "cryptonight", @@ -53,7 +55,9 @@ inline std::string get_algo_name(xmrstak_algo_id algo_id) "cryptonight_v8", "cryptonight_superfast", "cryptonight_gpu", - "cryptonight_conceal" + "cryptonight_conceal", + "cryptonight_r_wow", + "cryptonight_r" }}; static std::array derived_algo_names = @@ -174,7 +178,7 @@ constexpr uint32_t CN_ZELERIUS_ITER = 0x60000; inline xmrstak_algo POW(xmrstak_algo_id algo_id) { - static std::array pow = {{ + static std::array pow = {{ {invalid_algo, invalid_algo}, {cryptonight, cryptonight, CN_ITER, CN_MEMORY}, {cryptonight_lite, cryptonight_lite, CN_ITER/2, CN_MEMORY/2}, @@ -189,7 +193,9 @@ inline xmrstak_algo POW(xmrstak_algo_id algo_id) {cryptonight_monero_v8, cryptonight_monero_v8, CN_ITER, CN_MEMORY}, {cryptonight_superfast, cryptonight_superfast, CN_ITER/4, CN_MEMORY}, {cryptonight_gpu, cryptonight_gpu, CN_GPU_ITER, CN_MEMORY, CN_GPU_MASK}, - {cryptonight_conceal, cryptonight_conceal, CN_ITER/2, CN_MEMORY} + {cryptonight_conceal, cryptonight_conceal, CN_ITER/2, CN_MEMORY}, + {cryptonight_r_wow, cryptonight_r_wow, CN_ITER, CN_MEMORY}, + {cryptonight_r, cryptonight_r, CN_ITER, CN_MEMORY} }}; static std::array derived_pow = diff --git a/xmrstak/backend/globalStates.cpp b/xmrstak/backend/globalStates.cpp index 4eeed3c4b..52ef3f391 100644 --- a/xmrstak/backend/globalStates.cpp +++ b/xmrstak/backend/globalStates.cpp @@ -33,7 +33,7 @@ namespace xmrstak { -void globalStates::consume_work( miner_work& threadWork, uint64_t& currentJobId) +void globalStates::consume_work(miner_work& threadWork, uint64_t& currentJobId) { jobLock.ReadLock(); @@ -43,7 +43,7 @@ void globalStates::consume_work( miner_work& threadWork, uint64_t& currentJobId) jobLock.UnLock(); } -void globalStates::switch_work(miner_work& pWork, pool_data& dat) +void globalStates::switch_work(miner_work&& pWork, pool_data& dat) { jobLock.WriteLock(); @@ -61,7 +61,7 @@ void globalStates::switch_work(miner_work& pWork, pool_data& dat) * after the nonce is read. */ dat.iSavedNonce = iGlobalNonce.exchange(dat.iSavedNonce, std::memory_order_relaxed); - oGlobalWork = pWork; + oGlobalWork = std::move(pWork); jobLock.UnLock(); } diff --git a/xmrstak/backend/globalStates.hpp b/xmrstak/backend/globalStates.hpp index c8d691712..d6966c4a2 100644 --- a/xmrstak/backend/globalStates.hpp +++ b/xmrstak/backend/globalStates.hpp @@ -22,7 +22,7 @@ struct globalStates } //pool_data is in-out winapi style - void switch_work(miner_work& pWork, pool_data& dat); + void switch_work(miner_work&& pWork, pool_data& dat); inline void calc_start_nonce(uint32_t& nonce, bool use_nicehash, uint32_t reserve_count) { diff --git a/xmrstak/backend/miner_work.hpp b/xmrstak/backend/miner_work.hpp index b6456f031..c8174df32 100644 --- a/xmrstak/backend/miner_work.hpp +++ b/xmrstak/backend/miner_work.hpp @@ -21,29 +21,40 @@ namespace xmrstak bool bNiceHash; bool bStall; size_t iPoolId; + uint64_t iBlockHeight; + uint8_t* ref_ptr; - miner_work() : iWorkSize(0), bNiceHash(false), bStall(true), iPoolId(invalid_pool_id) { } + miner_work() : iWorkSize(0), bNiceHash(false), bStall(true), iPoolId(invalid_pool_id), ref_ptr((uint8_t*)&iBlockHeight) { } miner_work(const char* sJobID, const uint8_t* bWork, uint32_t iWorkSize, - uint64_t iTarget, bool bNiceHash, size_t iPoolId) : iWorkSize(iWorkSize), - iTarget(iTarget), bNiceHash(bNiceHash), bStall(false), iPoolId(iPoolId) + uint64_t iTarget, bool bNiceHash, size_t iPoolId, uint64_t iBlockHeiht) : iWorkSize(iWorkSize), + iTarget(iTarget), bNiceHash(bNiceHash), bStall(false), iPoolId(iPoolId), iBlockHeight(iBlockHeiht), ref_ptr((uint8_t*)&iBlockHeight) { assert(iWorkSize <= sizeof(bWorkBlob)); - memcpy(this->sJobID, sJobID, sizeof(miner_work::sJobID)); memcpy(this->bWorkBlob, bWork, iWorkSize); + memcpy(this->sJobID, sJobID, sizeof(miner_work::sJobID)); + } + + miner_work(miner_work&& from) : iWorkSize(from.iWorkSize), iTarget(from.iTarget), + bStall(from.bStall), iPoolId(from.iPoolId), iBlockHeight(from.iBlockHeight), ref_ptr((uint8_t*)&iBlockHeight) + { + assert(iWorkSize <= sizeof(bWorkBlob)); + memcpy(bWorkBlob, from.bWorkBlob, iWorkSize); + memcpy(this->sJobID, sJobID, sizeof(miner_work::sJobID)); } miner_work(miner_work const&) = delete; - miner_work& operator=(miner_work const& from) + miner_work& operator=(miner_work&& from) { assert(this != &from); + iBlockHeight = from.iBlockHeight; + iPoolId = from.iPoolId; + bStall = from.bStall; iWorkSize = from.iWorkSize; - iTarget = from.iTarget; bNiceHash = from.bNiceHash; - bStall = from.bStall; - iPoolId = from.iPoolId; + iTarget = from.iTarget; assert(iWorkSize <= sizeof(bWorkBlob)); memcpy(sJobID, from.sJobID, sizeof(sJobID)); @@ -52,23 +63,22 @@ namespace xmrstak return *this; } - miner_work(miner_work&& from) : iWorkSize(from.iWorkSize), iTarget(from.iTarget), - bStall(from.bStall), iPoolId(from.iPoolId) - { - assert(iWorkSize <= sizeof(bWorkBlob)); - memcpy(sJobID, from.sJobID, sizeof(sJobID)); - memcpy(bWorkBlob, from.bWorkBlob, iWorkSize); - } - - miner_work& operator=(miner_work&& from) + miner_work& operator=(miner_work const& from) { assert(this != &from); + iBlockHeight = from.iBlockHeight; + iPoolId = from.iPoolId; + bStall = from.bStall; iWorkSize = from.iWorkSize; - iTarget = from.iTarget; bNiceHash = from.bNiceHash; - bStall = from.bStall; - iPoolId = from.iPoolId; + iTarget = from.iTarget; + + if(!ref_ptr) + return *this; + + for(size_t i=0; i <= 7 && iPoolId; i++) + ref_ptr[i] = from.ref_ptr[7-i]; assert(iWorkSize <= sizeof(bWorkBlob)); memcpy(sJobID, from.sJobID, sizeof(sJobID)); diff --git a/xmrstak/backend/nvidia/CudaCryptonightR_gen.cpp b/xmrstak/backend/nvidia/CudaCryptonightR_gen.cpp new file mode 100644 index 000000000..87eb05540 --- /dev/null +++ b/xmrstak/backend/nvidia/CudaCryptonightR_gen.cpp @@ -0,0 +1,336 @@ +/* + * This program is free software: you can redistribute it and/or modify + * it under the terms of the GNU General Public License as published by + * the Free Software Foundation, either version 3 of the License, or + * any later version. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with this program. If not, see . + * + */ + +#include +#include +#include +#include +#include +#include + +#include "xmrstak/backend/nvidia/CudaCryptonightR_gen.hpp" +#include "xmrstak/backend/cpu/crypto/variant4_random_math.h" +#include "xmrstak/misc/console.hpp" +#include "xmrstak/cpputil/read_write_lock.h" + +namespace xmrstak +{ +namespace nvidia +{ + +static std::string get_code(const V4_Instruction* code, int code_size) +{ + std::stringstream s; + + for (int i = 0; i < code_size; ++i) + { + const V4_Instruction inst = code[i]; + + const uint32_t a = inst.dst_index; + const uint32_t b = inst.src_index; + + switch (inst.opcode) + { + case MUL: + s << 'r' << a << "*=r" << b << ';'; + break; + + case ADD: + s << 'r' << a << "+=r" << b << '+' << inst.C << "U;"; + break; + + case SUB: + s << 'r' << a << "-=r" << b << ';'; + break; + + case ROR: + s << 'r' << a << "=rotate_right(r" << a << ",r" << b << ");"; + break; + + case ROL: + s << 'r' << a << "=rotate_left(r" << a << ",r" << b << ");"; + break; + + case XOR: + s << 'r' << a << "^=r" << b << ';'; + break; + } + + s << '\n'; + } + + return s.str(); +} + +struct CacheEntry +{ + CacheEntry(xmrstak_algo algo, uint64_t height, int arch_major, int arch_minor, const std::vector& ptx, const std::string& lowered_name) : + algo(algo), + height(height), + arch_major(arch_major), + arch_minor(arch_minor), + ptx(ptx), + lowered_name(lowered_name) + {} + + xmrstak_algo algo; + uint64_t height; + int arch_major; + int arch_minor; + std::vector ptx; + std::string lowered_name; +}; + +struct BackgroundTaskBase +{ + virtual ~BackgroundTaskBase() {} + virtual void exec() = 0; +}; + +template +struct BackgroundTask : public BackgroundTaskBase +{ + BackgroundTask(T&& func) : m_func(std::move(func)) {} + void exec() override { m_func(); } + + T m_func; +}; + +static ::cpputil::RWLock CryptonightR_cache_mutex; +static std::mutex CryptonightR_build_mutex; +static std::vector CryptonightR_cache; + +static std::mutex background_tasks_mutex; +static std::vector background_tasks; +static std::thread* background_thread = nullptr; + +static void background_thread_proc() +{ + std::vector tasks; + for (;;) { + tasks.clear(); + { + std::lock_guard g(background_tasks_mutex); + background_tasks.swap(tasks); + } + + for (BackgroundTaskBase* task : tasks) { + task->exec(); + delete task; + } + + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + } +} + +template +static void background_exec(T&& func) +{ + BackgroundTaskBase* task = new BackgroundTask(std::move(func)); + + std::lock_guard g(background_tasks_mutex); + background_tasks.push_back(task); + if (!background_thread) { + background_thread = new std::thread(background_thread_proc); + } +} + +static void CryptonightR_build_program( + std::vector& ptx, + std::string& lowered_name, + const xmrstak_algo& algo, + uint64_t height, + int arch_major, + int arch_minor, + std::string source) +{ + { + CryptonightR_cache_mutex.WriteLock(); + + // Remove old programs from cache + for (size_t i = 0; i < CryptonightR_cache.size();) + { + const CacheEntry& entry = CryptonightR_cache[i]; + if ((entry.algo == algo) && (entry.height + 2 < height)) + { + printer::inst()->print_msg(LDEBUG, "CryptonightR: program for height %llu released (old program)", entry.height); + CryptonightR_cache[i] = std::move(CryptonightR_cache.back()); + CryptonightR_cache.pop_back(); + } + else + { + ++i; + } + } + CryptonightR_cache_mutex.UnLock(); + } + + ptx.clear(); + ptx.reserve(65536); + + std::lock_guard g1(CryptonightR_build_mutex); + { + CryptonightR_cache_mutex.ReadLock(); + + // Check if the cache already has this program (some other thread might have added it first) + for (const CacheEntry& entry : CryptonightR_cache) + { + if ((entry.algo == algo) && (entry.height == height) && (entry.arch_major == arch_major) && (entry.arch_minor == arch_minor)) + { + ptx = entry.ptx; + lowered_name = entry.lowered_name; + CryptonightR_cache_mutex.UnLock(); + return; + } + } + CryptonightR_cache_mutex.UnLock(); + } + + nvrtcProgram prog; + nvrtcResult result = nvrtcCreateProgram(&prog, source.c_str(), "CryptonightR.curt", 0, NULL, NULL); + if (result != NVRTC_SUCCESS) { + printer::inst()->print_msg(L0, "nvrtcCreateProgram failed: %s", nvrtcGetErrorString(result)); + return; + } + + result = nvrtcAddNameExpression(prog, "CryptonightR_phase2"); + if (result != NVRTC_SUCCESS) { + printer::inst()->print_msg(L0, "nvrtcAddNameExpression failed: %s", nvrtcGetErrorString(result)); + nvrtcDestroyProgram(&prog); + return; + } + + char opt0[64]; + sprintf(opt0, "--gpu-architecture=compute_%d%d", arch_major, arch_minor); + + char opt1[64]; + sprintf(opt1, "-DALGO=%d", static_cast(algo.Id())); + + const char* opts[2] = { opt0, opt1 }; + + result = nvrtcCompileProgram(prog, 2, opts); + if (result != NVRTC_SUCCESS) { + printer::inst()->print_msg(L0, "nvrtcCompileProgram failed: %s", nvrtcGetErrorString(result)); + + size_t logSize; + if (nvrtcGetProgramLogSize(prog, &logSize) == NVRTC_SUCCESS) { + char *log = new char[logSize]; + if (nvrtcGetProgramLog(prog, log) == NVRTC_SUCCESS) { + printer::inst()->print_msg(L0, "Program compile log: %s", log); + } + delete[]log; + } + nvrtcDestroyProgram(&prog); + return; + } + + + const char* name; + result = nvrtcGetLoweredName(prog, "CryptonightR_phase2", &name); + if (result != NVRTC_SUCCESS) { + printer::inst()->print_msg(L0, "nvrtcGetLoweredName failed: %s", nvrtcGetErrorString(result)); + nvrtcDestroyProgram(&prog); + return; + } + + size_t ptxSize; + result = nvrtcGetPTXSize(prog, &ptxSize); + if (result != NVRTC_SUCCESS) { + printer::inst()->print_msg(L0, "nvrtcGetPTXSize failed: %s", nvrtcGetErrorString(result)); + nvrtcDestroyProgram(&prog); + return; + } + + ptx.resize(ptxSize); + result = nvrtcGetPTX(prog, ptx.data()); + if (result != NVRTC_SUCCESS) { + printer::inst()->print_msg(L0, "nvrtcGetPTX failed: %s", nvrtcGetErrorString(result)); + nvrtcDestroyProgram(&prog); + return; + } + + lowered_name = name; + + nvrtcDestroyProgram(&prog); + + printer::inst()->print_msg(LDEBUG, "CryptonightR: program for height %llu compiled", height); + + CryptonightR_cache_mutex.WriteLock(); + CryptonightR_cache.emplace_back(algo, height, arch_major, arch_minor, ptx, lowered_name); + CryptonightR_cache_mutex.UnLock(); +} + +void CryptonightR_get_program(std::vector& ptx, std::string& lowered_name, const xmrstak_algo algo, uint64_t height, int arch_major, int arch_minor, bool background) +{ + if (background) { + background_exec([=]() { std::vector tmp; std::string s; CryptonightR_get_program(tmp, s, algo, height, arch_major, arch_minor, false); }); + return; + } + + ptx.clear(); + + const char* source_code_template = + #include "nvcc_code/cuda_cryptonight_r.curt" + ; + const char include_name[] = "XMRSTAK_INCLUDE_RANDOM_MATH"; + const char* offset = strstr(source_code_template, include_name); + if (!offset) + { + printer::inst()->print_msg(L0, "CryptonightR_get_program: XMRSTAK_INCLUDE_RANDOM_MATH not found in cuda_cryptonight_r.curt"); + return; + } + + V4_Instruction code[256]; + int code_size; + switch (algo.Id()) + { + case cryptonight_r_wow: + code_size = v4_random_math_init(code, height); + break; + case cryptonight_r: + code_size = v4_random_math_init(code, height); + break; + printer::inst()->print_msg(LDEBUG, "CryptonightR_get_program: invalid algo %d", algo); + return; + } + + std::string source_code(source_code_template, offset); + source_code.append(get_code(code, code_size)); + source_code.append(offset + sizeof(include_name) - 1); + + { + CryptonightR_cache_mutex.ReadLock(); + + // Check if the cache has this program + for (const CacheEntry& entry : CryptonightR_cache) + { + if ((entry.algo == algo) && (entry.height == height) && (entry.arch_major == arch_major) && (entry.arch_minor == arch_minor)) + { + printer::inst()->print_msg(LDEBUG, "CryptonightR: program for height %llu found in cache", height); + ptx = entry.ptx; + lowered_name = entry.lowered_name; + CryptonightR_cache_mutex.UnLock(); + return; + } + } + CryptonightR_cache_mutex.UnLock(); + } + + CryptonightR_build_program(ptx, lowered_name, algo, height, arch_major, arch_minor, source_code); +} + +} // namespace xmrstak +} //namespace nvidia diff --git a/xmrstak/backend/nvidia/CudaCryptonightR_gen.hpp b/xmrstak/backend/nvidia/CudaCryptonightR_gen.hpp new file mode 100644 index 000000000..e214647b9 --- /dev/null +++ b/xmrstak/backend/nvidia/CudaCryptonightR_gen.hpp @@ -0,0 +1,37 @@ +/* + * This program is free software: you can redistribute it and/or modify + * it under the terms of the GNU General Public License as published by + * the Free Software Foundation, either version 3 of the License, or + * any later version. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with this program. If not, see . + * + */ + +#pragma once + +#include "xmrstak/backend/cryptonight.hpp" + +#include +#include +#include + + +namespace xmrstak +{ +namespace nvidia +{ + +void CryptonightR_get_program(std::vector& ptx, std::string& lowered_name, + const xmrstak_algo algo, uint64_t height, int arch_major, int arch_minor, bool background = false); + + +} // namespace xmrstak +} //namespace nvidia + diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp index 07ed4d31e..794e68d11 100644 --- a/xmrstak/backend/nvidia/minethd.cpp +++ b/xmrstak/backend/nvidia/minethd.cpp @@ -173,6 +173,8 @@ std::vector* minethd::thread_starter(uint32_t threadOffset, miner_wor size_t i, n = jconf::inst()->GetGPUThreadCount(); pvThreads->reserve(n); + cuInit(0); + jconf::thd_cfg cfg; for (i = 0; i < n; i++) { @@ -226,7 +228,10 @@ void minethd::work_main() // start with root algorithm and switch later if fork version is reached 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); + + cn_hash_fun hash_fun; + cpu::minethd::cn_on_new_job set_job; + cpu::minethd::func_multi_selector<1>(hash_fun, set_job, ::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); uint32_t iNonce; @@ -255,17 +260,20 @@ void minethd::work_main() if(new_version >= coinDesc.GetMiningForkVersion()) { miner_algo = coinDesc.GetMiningAlgo(); - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); + cpu::minethd::func_multi_selector<1>(hash_fun, set_job, ::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); } else { miner_algo = coinDesc.GetMiningAlgoRoot(); - hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); + cpu::minethd::func_multi_selector<1>(hash_fun, set_job, ::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo); } lastPoolId = oWork.iPoolId; version = new_version; } + if(set_job != nullptr) + set_job(oWork, &cpu_ctx); + cryptonight_extra_cpu_set_data(&ctx, oWork.bWorkBlob, oWork.iWorkSize); uint32_t h_per_round = ctx.device_blocks * ctx.device_threads; @@ -292,7 +300,7 @@ void minethd::work_main() cryptonight_extra_cpu_prepare(&ctx, iNonce, miner_algo); - cryptonight_core_cpu_hash(&ctx, miner_algo, iNonce); + cryptonight_core_cpu_hash(&ctx, miner_algo, iNonce, cpu_ctx->cn_r_ctx.height); cryptonight_extra_cpu_final(&ctx, iNonce, oWork.iTarget, &foundCount, foundNonce, miner_algo); diff --git a/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp b/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp index 45ffef806..fe77b6f81 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp +++ b/xmrstak/backend/nvidia/nvcc_code/cryptonight.hpp @@ -6,6 +6,8 @@ #include "xmrstak/jconf.hpp" #include "xmrstak/backend/cryptonight.hpp" +#include + typedef struct { int device_id; const char *device_name; @@ -33,6 +35,13 @@ typedef struct { std::string name; size_t free_device_memory; size_t total_device_memory; + + CUdevice cuDevice; + CUcontext cuContext; + CUmodule module = nullptr; + CUfunction kernel = nullptr; + uint64_t kernel_height = 0; + xmrstak_algo cached_algo = {xmrstak_algo_id::invalid_algo}; } nvid_ctx; extern "C" { @@ -50,4 +59,4 @@ void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce, const xmr void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce, const xmrstak_algo& miner_algo); } -void cryptonight_core_cpu_hash(nvid_ctx* ctx, const xmrstak_algo& miner_algo, uint32_t startNonce); +void cryptonight_core_cpu_hash(nvid_ctx* ctx, const xmrstak_algo& miner_algo, uint32_t startNonce, uint64_t chain_height); diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu index 184825222..d082f3362 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu @@ -11,6 +11,7 @@ #include "xmrstak/backend/nvidia/nvcc_code/cuda_fast_int_math_v2.hpp" #include "xmrstak/backend/nvidia/nvcc_code/cuda_fast_div_heavy.hpp" #include "xmrstak/backend/nvidia/nvcc_code/cuda_cryptonight_gpu.hpp" +#include "xmrstak/backend/nvidia/CudaCryptonightR_gen.hpp" #ifdef _WIN32 @@ -316,11 +317,11 @@ __global__ void cryptonight_core_gpu_phase2_double( uint64_t division_result; if(ALGO == cryptonight_monero_v8) { - bx0 = ((uint64_t*)(d_ctx_b + thread * 12))[sub]; - bx1 = ((uint64_t*)(d_ctx_b + thread * 12 + 4))[sub]; + bx0 = ((uint64_t*)(d_ctx_b + thread * 16))[sub]; + bx1 = ((uint64_t*)(d_ctx_b + thread * 16 + 4))[sub]; - division_result = ((uint64_t*)(d_ctx_b + thread * 12 + 4 * 2))[0]; - sqrt_result = (d_ctx_b + thread * 12 + 4 * 2 + 2)[0]; + division_result = ((uint64_t*)(d_ctx_b + thread * 16 + 4 * 2))[0]; + sqrt_result = (d_ctx_b + thread * 16 + 4 * 2 + 2)[0]; } else bx0 = ((uint64_t*)(d_ctx_b + thread * 4))[sub]; @@ -470,14 +471,14 @@ __global__ void cryptonight_core_gpu_phase2_double( ((uint64_t*)(d_ctx_a + thread * 4))[sub] = ax0; if(ALGO == cryptonight_monero_v8) { - ((uint64_t*)(d_ctx_b + thread * 12))[sub] = bx0; - ((uint64_t*)(d_ctx_b + thread * 12 + 4))[sub] = bx1; + ((uint64_t*)(d_ctx_b + thread * 16))[sub] = bx0; + ((uint64_t*)(d_ctx_b + thread * 16 + 4))[sub] = bx1; if(sub == 1) { // must be valid only for `sub == 1` - ((uint64_t*)(d_ctx_b + thread * 12 + 4 * 2))[0] = division_result; - (d_ctx_b + thread * 12 + 4 * 2 + 2)[0] = sqrt_result; + ((uint64_t*)(d_ctx_b + thread * 16 + 4 * 2))[0] = division_result; + (d_ctx_b + thread * 16 + 4 * 2 + 2)[0] = sqrt_result; } } else @@ -531,7 +532,7 @@ __global__ void cryptonight_core_gpu_phase2_quad( else conc_var = 0.0f; } - + uint32_t tweak1_2[2]; if (ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_masari || ALGO == cryptonight_bittube2) { @@ -600,7 +601,7 @@ __global__ void cryptonight_core_gpu_phase2_quad( { float r = int2float((int32_t)x_0); float c_old = conc_var; - + r += conc_var; r = r * r * r; r = int_as_float((float_as_int(r) & 0x807FFFFF) | 0x40000000); @@ -774,9 +775,9 @@ __global__ void cryptonight_core_gpu_phase3( template void cryptonight_core_gpu_hash(nvid_ctx* ctx, uint32_t nonce, const xmrstak_algo& algo) { - const uint32_t MASK = algo.Mask(); - const uint32_t ITERATIONS = algo.Iter(); - const size_t MEM = algo.Mem()/4; + uint32_t MASK = algo.Mask(); + uint32_t ITERATIONS = algo.Iter(); + size_t MEM = algo.Mem()/4; dim3 grid( ctx->device_blocks ); dim3 block( ctx->device_threads ); @@ -823,7 +824,7 @@ void cryptonight_core_gpu_hash(nvid_ctx* ctx, uint32_t nonce, const xmrstak_algo cryptonight_core_gpu_phase2_double<<< grid, block2, - sizeof(uint64_t) * block2.x * 8 + + sizeof(uint64_t) * block.x * 8 + // shuffle memory for fermi gpus block2.x * sizeof(uint32_t) * static_cast< int >( ctx->device_arch[0] < 3 ) >>>( @@ -842,6 +843,26 @@ void cryptonight_core_gpu_hash(nvid_ctx* ctx, uint32_t nonce, const xmrstak_algo ) ); } + else if(ALGO == cryptonight_r_wow || ALGO == cryptonight_r) + { + int numThreads = ctx->device_blocks*ctx->device_threads; + void* args[] = { + &ITERATIONS, &MEM, &MASK, + &numThreads, &ctx->device_bfactor, &i, + &ctx->d_long_state, &ctx->d_ctx_a, &ctx->d_ctx_b, &ctx->d_ctx_state, &nonce, &ctx->d_input + }; + CU_CHECK(ctx->device_id, cuLaunchKernel( + ctx->kernel, + grid.x, grid.y, grid.z, + block2.x, block2.y, block2.z, + sizeof(uint64_t) * block.x * 8 + + // shuffle memory for fermi gpus + block2.x * sizeof(uint32_t) * static_cast< int >( ctx->device_arch[0] < 3 ), + nullptr, + args, 0 + )); + CU_CHECK(ctx->device_id, cuCtxSynchronize()); + } else { CUDA_CHECK_MSG_KERNEL( @@ -972,8 +993,30 @@ void cryptonight_core_gpu_hash_gpu(nvid_ctx* ctx, uint32_t nonce, const xmrstak_ } } -void cryptonight_core_cpu_hash(nvid_ctx* ctx, const xmrstak_algo& miner_algo, uint32_t startNonce) +void cryptonight_core_cpu_hash(nvid_ctx* ctx, const xmrstak_algo& miner_algo, uint32_t startNonce, uint64_t chain_height) { + + if((miner_algo == cryptonight_r_wow) || (miner_algo == cryptonight_r)) + { + if(ctx->kernel_height != chain_height || ctx->cached_algo != miner_algo) + { + if(ctx->module) + cuModuleUnload(ctx->module); + + std::vector ptx; + std::string lowered_name; + xmrstak::nvidia::CryptonightR_get_program(ptx, lowered_name, miner_algo, chain_height, ctx->device_arch[0], ctx->device_arch[1]); + + CU_CHECK(ctx->device_id, cuModuleLoadDataEx(&ctx->module, ptx.data(), 0, 0, 0)); + CU_CHECK(ctx->device_id, cuModuleGetFunction(&ctx->kernel, ctx->module, lowered_name.c_str())); + + ctx->kernel_height = chain_height; + ctx->cached_algo = miner_algo; + + xmrstak::nvidia::CryptonightR_get_program(ptx, lowered_name, miner_algo, chain_height + 1, ctx->device_arch[0], ctx->device_arch[1], true); + } + } + typedef void (*cuda_hash_fn)(nvid_ctx* ctx, uint32_t nonce, const xmrstak_algo& algo); if(miner_algo == invalid_algo) return; @@ -1019,7 +1062,13 @@ void cryptonight_core_cpu_hash(nvid_ctx* ctx, const xmrstak_algo& miner_algo, ui cryptonight_core_gpu_hash_gpu, cryptonight_core_gpu_hash, - cryptonight_core_gpu_hash + cryptonight_core_gpu_hash, + + cryptonight_core_gpu_hash, + cryptonight_core_gpu_hash, + + cryptonight_core_gpu_hash, + cryptonight_core_gpu_hash }; std::bitset<1> digit; diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_cryptonight_r.curt b/xmrstak/backend/nvidia/nvcc_code/cuda_cryptonight_r.curt new file mode 100644 index 000000000..bcf495080 --- /dev/null +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_cryptonight_r.curt @@ -0,0 +1,618 @@ +R"===( + +/* + * This program is free software: you can redistribute it and/or modify + * it under the terms of the GNU General Public License as published by + * the Free Software Foundation, either version 3 of the License, or + * any later version. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with this program. If not, see . + * + */ + +#define cryptonight_r_wow 15 +#define cryptonight_r 16 + +typedef unsigned char uint8_t; +typedef unsigned int uint32_t; +typedef unsigned long long int uint64_t; + +static __constant__ uint32_t d_t_fn[1024] = +{ + 0xa56363c6U, 0x847c7cf8U, 0x997777eeU, 0x8d7b7bf6U, + 0x0df2f2ffU, 0xbd6b6bd6U, 0xb16f6fdeU, 0x54c5c591U, + 0x50303060U, 0x03010102U, 0xa96767ceU, 0x7d2b2b56U, + 0x19fefee7U, 0x62d7d7b5U, 0xe6abab4dU, 0x9a7676ecU, + 0x45caca8fU, 0x9d82821fU, 0x40c9c989U, 0x877d7dfaU, + 0x15fafaefU, 0xeb5959b2U, 0xc947478eU, 0x0bf0f0fbU, + 0xecadad41U, 0x67d4d4b3U, 0xfda2a25fU, 0xeaafaf45U, + 0xbf9c9c23U, 0xf7a4a453U, 0x967272e4U, 0x5bc0c09bU, + 0xc2b7b775U, 0x1cfdfde1U, 0xae93933dU, 0x6a26264cU, + 0x5a36366cU, 0x413f3f7eU, 0x02f7f7f5U, 0x4fcccc83U, + 0x5c343468U, 0xf4a5a551U, 0x34e5e5d1U, 0x08f1f1f9U, + 0x937171e2U, 0x73d8d8abU, 0x53313162U, 0x3f15152aU, + 0x0c040408U, 0x52c7c795U, 0x65232346U, 0x5ec3c39dU, + 0x28181830U, 0xa1969637U, 0x0f05050aU, 0xb59a9a2fU, + 0x0907070eU, 0x36121224U, 0x9b80801bU, 0x3de2e2dfU, + 0x26ebebcdU, 0x6927274eU, 0xcdb2b27fU, 0x9f7575eaU, + 0x1b090912U, 0x9e83831dU, 0x742c2c58U, 0x2e1a1a34U, + 0x2d1b1b36U, 0xb26e6edcU, 0xee5a5ab4U, 0xfba0a05bU, + 0xf65252a4U, 0x4d3b3b76U, 0x61d6d6b7U, 0xceb3b37dU, + 0x7b292952U, 0x3ee3e3ddU, 0x712f2f5eU, 0x97848413U, + 0xf55353a6U, 0x68d1d1b9U, 0x00000000U, 0x2cededc1U, + 0x60202040U, 0x1ffcfce3U, 0xc8b1b179U, 0xed5b5bb6U, + 0xbe6a6ad4U, 0x46cbcb8dU, 0xd9bebe67U, 0x4b393972U, + 0xde4a4a94U, 0xd44c4c98U, 0xe85858b0U, 0x4acfcf85U, + 0x6bd0d0bbU, 0x2aefefc5U, 0xe5aaaa4fU, 0x16fbfbedU, + 0xc5434386U, 0xd74d4d9aU, 0x55333366U, 0x94858511U, + 0xcf45458aU, 0x10f9f9e9U, 0x06020204U, 0x817f7ffeU, + 0xf05050a0U, 0x443c3c78U, 0xba9f9f25U, 0xe3a8a84bU, + 0xf35151a2U, 0xfea3a35dU, 0xc0404080U, 0x8a8f8f05U, + 0xad92923fU, 0xbc9d9d21U, 0x48383870U, 0x04f5f5f1U, + 0xdfbcbc63U, 0xc1b6b677U, 0x75dadaafU, 0x63212142U, + 0x30101020U, 0x1affffe5U, 0x0ef3f3fdU, 0x6dd2d2bfU, + 0x4ccdcd81U, 0x140c0c18U, 0x35131326U, 0x2fececc3U, + 0xe15f5fbeU, 0xa2979735U, 0xcc444488U, 0x3917172eU, + 0x57c4c493U, 0xf2a7a755U, 0x827e7efcU, 0x473d3d7aU, + 0xac6464c8U, 0xe75d5dbaU, 0x2b191932U, 0x957373e6U, + 0xa06060c0U, 0x98818119U, 0xd14f4f9eU, 0x7fdcdca3U, + 0x66222244U, 0x7e2a2a54U, 0xab90903bU, 0x8388880bU, + 0xca46468cU, 0x29eeeec7U, 0xd3b8b86bU, 0x3c141428U, + 0x79dedea7U, 0xe25e5ebcU, 0x1d0b0b16U, 0x76dbdbadU, + 0x3be0e0dbU, 0x56323264U, 0x4e3a3a74U, 0x1e0a0a14U, + 0xdb494992U, 0x0a06060cU, 0x6c242448U, 0xe45c5cb8U, + 0x5dc2c29fU, 0x6ed3d3bdU, 0xefacac43U, 0xa66262c4U, + 0xa8919139U, 0xa4959531U, 0x37e4e4d3U, 0x8b7979f2U, + 0x32e7e7d5U, 0x43c8c88bU, 0x5937376eU, 0xb76d6ddaU, + 0x8c8d8d01U, 0x64d5d5b1U, 0xd24e4e9cU, 0xe0a9a949U, + 0xb46c6cd8U, 0xfa5656acU, 0x07f4f4f3U, 0x25eaeacfU, + 0xaf6565caU, 0x8e7a7af4U, 0xe9aeae47U, 0x18080810U, + 0xd5baba6fU, 0x887878f0U, 0x6f25254aU, 0x722e2e5cU, + 0x241c1c38U, 0xf1a6a657U, 0xc7b4b473U, 0x51c6c697U, + 0x23e8e8cbU, 0x7cdddda1U, 0x9c7474e8U, 0x211f1f3eU, + 0xdd4b4b96U, 0xdcbdbd61U, 0x868b8b0dU, 0x858a8a0fU, + 0x907070e0U, 0x423e3e7cU, 0xc4b5b571U, 0xaa6666ccU, + 0xd8484890U, 0x05030306U, 0x01f6f6f7U, 0x120e0e1cU, + 0xa36161c2U, 0x5f35356aU, 0xf95757aeU, 0xd0b9b969U, + 0x91868617U, 0x58c1c199U, 0x271d1d3aU, 0xb99e9e27U, + 0x38e1e1d9U, 0x13f8f8ebU, 0xb398982bU, 0x33111122U, + 0xbb6969d2U, 0x70d9d9a9U, 0x898e8e07U, 0xa7949433U, + 0xb69b9b2dU, 0x221e1e3cU, 0x92878715U, 0x20e9e9c9U, + 0x49cece87U, 0xff5555aaU, 0x78282850U, 0x7adfdfa5U, + 0x8f8c8c03U, 0xf8a1a159U, 0x80898909U, 0x170d0d1aU, + 0xdabfbf65U, 0x31e6e6d7U, 0xc6424284U, 0xb86868d0U, + 0xc3414182U, 0xb0999929U, 0x772d2d5aU, 0x110f0f1eU, + 0xcbb0b07bU, 0xfc5454a8U, 0xd6bbbb6dU, 0x3a16162cU, + 0x6363c6a5U, 0x7c7cf884U, 0x7777ee99U, 0x7b7bf68dU, + 0xf2f2ff0dU, 0x6b6bd6bdU, 0x6f6fdeb1U, 0xc5c59154U, + 0x30306050U, 0x01010203U, 0x6767cea9U, 0x2b2b567dU, + 0xfefee719U, 0xd7d7b562U, 0xabab4de6U, 0x7676ec9aU, + 0xcaca8f45U, 0x82821f9dU, 0xc9c98940U, 0x7d7dfa87U, + 0xfafaef15U, 0x5959b2ebU, 0x47478ec9U, 0xf0f0fb0bU, + 0xadad41ecU, 0xd4d4b367U, 0xa2a25ffdU, 0xafaf45eaU, + 0x9c9c23bfU, 0xa4a453f7U, 0x7272e496U, 0xc0c09b5bU, + 0xb7b775c2U, 0xfdfde11cU, 0x93933daeU, 0x26264c6aU, + 0x36366c5aU, 0x3f3f7e41U, 0xf7f7f502U, 0xcccc834fU, + 0x3434685cU, 0xa5a551f4U, 0xe5e5d134U, 0xf1f1f908U, + 0x7171e293U, 0xd8d8ab73U, 0x31316253U, 0x15152a3fU, + 0x0404080cU, 0xc7c79552U, 0x23234665U, 0xc3c39d5eU, + 0x18183028U, 0x969637a1U, 0x05050a0fU, 0x9a9a2fb5U, + 0x07070e09U, 0x12122436U, 0x80801b9bU, 0xe2e2df3dU, + 0xebebcd26U, 0x27274e69U, 0xb2b27fcdU, 0x7575ea9fU, + 0x0909121bU, 0x83831d9eU, 0x2c2c5874U, 0x1a1a342eU, + 0x1b1b362dU, 0x6e6edcb2U, 0x5a5ab4eeU, 0xa0a05bfbU, + 0x5252a4f6U, 0x3b3b764dU, 0xd6d6b761U, 0xb3b37dceU, + 0x2929527bU, 0xe3e3dd3eU, 0x2f2f5e71U, 0x84841397U, + 0x5353a6f5U, 0xd1d1b968U, 0x00000000U, 0xededc12cU, + 0x20204060U, 0xfcfce31fU, 0xb1b179c8U, 0x5b5bb6edU, + 0x6a6ad4beU, 0xcbcb8d46U, 0xbebe67d9U, 0x3939724bU, + 0x4a4a94deU, 0x4c4c98d4U, 0x5858b0e8U, 0xcfcf854aU, + 0xd0d0bb6bU, 0xefefc52aU, 0xaaaa4fe5U, 0xfbfbed16U, + 0x434386c5U, 0x4d4d9ad7U, 0x33336655U, 0x85851194U, + 0x45458acfU, 0xf9f9e910U, 0x02020406U, 0x7f7ffe81U, + 0x5050a0f0U, 0x3c3c7844U, 0x9f9f25baU, 0xa8a84be3U, + 0x5151a2f3U, 0xa3a35dfeU, 0x404080c0U, 0x8f8f058aU, + 0x92923fadU, 0x9d9d21bcU, 0x38387048U, 0xf5f5f104U, + 0xbcbc63dfU, 0xb6b677c1U, 0xdadaaf75U, 0x21214263U, + 0x10102030U, 0xffffe51aU, 0xf3f3fd0eU, 0xd2d2bf6dU, + 0xcdcd814cU, 0x0c0c1814U, 0x13132635U, 0xececc32fU, + 0x5f5fbee1U, 0x979735a2U, 0x444488ccU, 0x17172e39U, + 0xc4c49357U, 0xa7a755f2U, 0x7e7efc82U, 0x3d3d7a47U, + 0x6464c8acU, 0x5d5dbae7U, 0x1919322bU, 0x7373e695U, + 0x6060c0a0U, 0x81811998U, 0x4f4f9ed1U, 0xdcdca37fU, + 0x22224466U, 0x2a2a547eU, 0x90903babU, 0x88880b83U, + 0x46468ccaU, 0xeeeec729U, 0xb8b86bd3U, 0x1414283cU, + 0xdedea779U, 0x5e5ebce2U, 0x0b0b161dU, 0xdbdbad76U, + 0xe0e0db3bU, 0x32326456U, 0x3a3a744eU, 0x0a0a141eU, + 0x494992dbU, 0x06060c0aU, 0x2424486cU, 0x5c5cb8e4U, + 0xc2c29f5dU, 0xd3d3bd6eU, 0xacac43efU, 0x6262c4a6U, + 0x919139a8U, 0x959531a4U, 0xe4e4d337U, 0x7979f28bU, + 0xe7e7d532U, 0xc8c88b43U, 0x37376e59U, 0x6d6ddab7U, + 0x8d8d018cU, 0xd5d5b164U, 0x4e4e9cd2U, 0xa9a949e0U, + 0x6c6cd8b4U, 0x5656acfaU, 0xf4f4f307U, 0xeaeacf25U, + 0x6565caafU, 0x7a7af48eU, 0xaeae47e9U, 0x08081018U, + 0xbaba6fd5U, 0x7878f088U, 0x25254a6fU, 0x2e2e5c72U, + 0x1c1c3824U, 0xa6a657f1U, 0xb4b473c7U, 0xc6c69751U, + 0xe8e8cb23U, 0xdddda17cU, 0x7474e89cU, 0x1f1f3e21U, + 0x4b4b96ddU, 0xbdbd61dcU, 0x8b8b0d86U, 0x8a8a0f85U, + 0x7070e090U, 0x3e3e7c42U, 0xb5b571c4U, 0x6666ccaaU, + 0x484890d8U, 0x03030605U, 0xf6f6f701U, 0x0e0e1c12U, + 0x6161c2a3U, 0x35356a5fU, 0x5757aef9U, 0xb9b969d0U, + 0x86861791U, 0xc1c19958U, 0x1d1d3a27U, 0x9e9e27b9U, + 0xe1e1d938U, 0xf8f8eb13U, 0x98982bb3U, 0x11112233U, + 0x6969d2bbU, 0xd9d9a970U, 0x8e8e0789U, 0x949433a7U, + 0x9b9b2db6U, 0x1e1e3c22U, 0x87871592U, 0xe9e9c920U, + 0xcece8749U, 0x5555aaffU, 0x28285078U, 0xdfdfa57aU, + 0x8c8c038fU, 0xa1a159f8U, 0x89890980U, 0x0d0d1a17U, + 0xbfbf65daU, 0xe6e6d731U, 0x424284c6U, 0x6868d0b8U, + 0x414182c3U, 0x999929b0U, 0x2d2d5a77U, 0x0f0f1e11U, + 0xb0b07bcbU, 0x5454a8fcU, 0xbbbb6dd6U, 0x16162c3aU, + 0x63c6a563U, 0x7cf8847cU, 0x77ee9977U, 0x7bf68d7bU, + 0xf2ff0df2U, 0x6bd6bd6bU, 0x6fdeb16fU, 0xc59154c5U, + 0x30605030U, 0x01020301U, 0x67cea967U, 0x2b567d2bU, + 0xfee719feU, 0xd7b562d7U, 0xab4de6abU, 0x76ec9a76U, + 0xca8f45caU, 0x821f9d82U, 0xc98940c9U, 0x7dfa877dU, + 0xfaef15faU, 0x59b2eb59U, 0x478ec947U, 0xf0fb0bf0U, + 0xad41ecadU, 0xd4b367d4U, 0xa25ffda2U, 0xaf45eaafU, + 0x9c23bf9cU, 0xa453f7a4U, 0x72e49672U, 0xc09b5bc0U, + 0xb775c2b7U, 0xfde11cfdU, 0x933dae93U, 0x264c6a26U, + 0x366c5a36U, 0x3f7e413fU, 0xf7f502f7U, 0xcc834fccU, + 0x34685c34U, 0xa551f4a5U, 0xe5d134e5U, 0xf1f908f1U, + 0x71e29371U, 0xd8ab73d8U, 0x31625331U, 0x152a3f15U, + 0x04080c04U, 0xc79552c7U, 0x23466523U, 0xc39d5ec3U, + 0x18302818U, 0x9637a196U, 0x050a0f05U, 0x9a2fb59aU, + 0x070e0907U, 0x12243612U, 0x801b9b80U, 0xe2df3de2U, + 0xebcd26ebU, 0x274e6927U, 0xb27fcdb2U, 0x75ea9f75U, + 0x09121b09U, 0x831d9e83U, 0x2c58742cU, 0x1a342e1aU, + 0x1b362d1bU, 0x6edcb26eU, 0x5ab4ee5aU, 0xa05bfba0U, + 0x52a4f652U, 0x3b764d3bU, 0xd6b761d6U, 0xb37dceb3U, + 0x29527b29U, 0xe3dd3ee3U, 0x2f5e712fU, 0x84139784U, + 0x53a6f553U, 0xd1b968d1U, 0x00000000U, 0xedc12cedU, + 0x20406020U, 0xfce31ffcU, 0xb179c8b1U, 0x5bb6ed5bU, + 0x6ad4be6aU, 0xcb8d46cbU, 0xbe67d9beU, 0x39724b39U, + 0x4a94de4aU, 0x4c98d44cU, 0x58b0e858U, 0xcf854acfU, + 0xd0bb6bd0U, 0xefc52aefU, 0xaa4fe5aaU, 0xfbed16fbU, + 0x4386c543U, 0x4d9ad74dU, 0x33665533U, 0x85119485U, + 0x458acf45U, 0xf9e910f9U, 0x02040602U, 0x7ffe817fU, + 0x50a0f050U, 0x3c78443cU, 0x9f25ba9fU, 0xa84be3a8U, + 0x51a2f351U, 0xa35dfea3U, 0x4080c040U, 0x8f058a8fU, + 0x923fad92U, 0x9d21bc9dU, 0x38704838U, 0xf5f104f5U, + 0xbc63dfbcU, 0xb677c1b6U, 0xdaaf75daU, 0x21426321U, + 0x10203010U, 0xffe51affU, 0xf3fd0ef3U, 0xd2bf6dd2U, + 0xcd814ccdU, 0x0c18140cU, 0x13263513U, 0xecc32fecU, + 0x5fbee15fU, 0x9735a297U, 0x4488cc44U, 0x172e3917U, + 0xc49357c4U, 0xa755f2a7U, 0x7efc827eU, 0x3d7a473dU, + 0x64c8ac64U, 0x5dbae75dU, 0x19322b19U, 0x73e69573U, + 0x60c0a060U, 0x81199881U, 0x4f9ed14fU, 0xdca37fdcU, + 0x22446622U, 0x2a547e2aU, 0x903bab90U, 0x880b8388U, + 0x468cca46U, 0xeec729eeU, 0xb86bd3b8U, 0x14283c14U, + 0xdea779deU, 0x5ebce25eU, 0x0b161d0bU, 0xdbad76dbU, + 0xe0db3be0U, 0x32645632U, 0x3a744e3aU, 0x0a141e0aU, + 0x4992db49U, 0x060c0a06U, 0x24486c24U, 0x5cb8e45cU, + 0xc29f5dc2U, 0xd3bd6ed3U, 0xac43efacU, 0x62c4a662U, + 0x9139a891U, 0x9531a495U, 0xe4d337e4U, 0x79f28b79U, + 0xe7d532e7U, 0xc88b43c8U, 0x376e5937U, 0x6ddab76dU, + 0x8d018c8dU, 0xd5b164d5U, 0x4e9cd24eU, 0xa949e0a9U, + 0x6cd8b46cU, 0x56acfa56U, 0xf4f307f4U, 0xeacf25eaU, + 0x65caaf65U, 0x7af48e7aU, 0xae47e9aeU, 0x08101808U, + 0xba6fd5baU, 0x78f08878U, 0x254a6f25U, 0x2e5c722eU, + 0x1c38241cU, 0xa657f1a6U, 0xb473c7b4U, 0xc69751c6U, + 0xe8cb23e8U, 0xdda17cddU, 0x74e89c74U, 0x1f3e211fU, + 0x4b96dd4bU, 0xbd61dcbdU, 0x8b0d868bU, 0x8a0f858aU, + 0x70e09070U, 0x3e7c423eU, 0xb571c4b5U, 0x66ccaa66U, + 0x4890d848U, 0x03060503U, 0xf6f701f6U, 0x0e1c120eU, + 0x61c2a361U, 0x356a5f35U, 0x57aef957U, 0xb969d0b9U, + 0x86179186U, 0xc19958c1U, 0x1d3a271dU, 0x9e27b99eU, + 0xe1d938e1U, 0xf8eb13f8U, 0x982bb398U, 0x11223311U, + 0x69d2bb69U, 0xd9a970d9U, 0x8e07898eU, 0x9433a794U, + 0x9b2db69bU, 0x1e3c221eU, 0x87159287U, 0xe9c920e9U, + 0xce8749ceU, 0x55aaff55U, 0x28507828U, 0xdfa57adfU, + 0x8c038f8cU, 0xa159f8a1U, 0x89098089U, 0x0d1a170dU, + 0xbf65dabfU, 0xe6d731e6U, 0x4284c642U, 0x68d0b868U, + 0x4182c341U, 0x9929b099U, 0x2d5a772dU, 0x0f1e110fU, + 0xb07bcbb0U, 0x54a8fc54U, 0xbb6dd6bbU, 0x162c3a16U, + 0xc6a56363U, 0xf8847c7cU, 0xee997777U, 0xf68d7b7bU, + 0xff0df2f2U, 0xd6bd6b6bU, 0xdeb16f6fU, 0x9154c5c5U, + 0x60503030U, 0x02030101U, 0xcea96767U, 0x567d2b2bU, + 0xe719fefeU, 0xb562d7d7U, 0x4de6ababU, 0xec9a7676U, + 0x8f45cacaU, 0x1f9d8282U, 0x8940c9c9U, 0xfa877d7dU, + 0xef15fafaU, 0xb2eb5959U, 0x8ec94747U, 0xfb0bf0f0U, + 0x41ecadadU, 0xb367d4d4U, 0x5ffda2a2U, 0x45eaafafU, + 0x23bf9c9cU, 0x53f7a4a4U, 0xe4967272U, 0x9b5bc0c0U, + 0x75c2b7b7U, 0xe11cfdfdU, 0x3dae9393U, 0x4c6a2626U, + 0x6c5a3636U, 0x7e413f3fU, 0xf502f7f7U, 0x834fccccU, + 0x685c3434U, 0x51f4a5a5U, 0xd134e5e5U, 0xf908f1f1U, + 0xe2937171U, 0xab73d8d8U, 0x62533131U, 0x2a3f1515U, + 0x080c0404U, 0x9552c7c7U, 0x46652323U, 0x9d5ec3c3U, + 0x30281818U, 0x37a19696U, 0x0a0f0505U, 0x2fb59a9aU, + 0x0e090707U, 0x24361212U, 0x1b9b8080U, 0xdf3de2e2U, + 0xcd26ebebU, 0x4e692727U, 0x7fcdb2b2U, 0xea9f7575U, + 0x121b0909U, 0x1d9e8383U, 0x58742c2cU, 0x342e1a1aU, + 0x362d1b1bU, 0xdcb26e6eU, 0xb4ee5a5aU, 0x5bfba0a0U, + 0xa4f65252U, 0x764d3b3bU, 0xb761d6d6U, 0x7dceb3b3U, + 0x527b2929U, 0xdd3ee3e3U, 0x5e712f2fU, 0x13978484U, + 0xa6f55353U, 0xb968d1d1U, 0x00000000U, 0xc12cededU, + 0x40602020U, 0xe31ffcfcU, 0x79c8b1b1U, 0xb6ed5b5bU, + 0xd4be6a6aU, 0x8d46cbcbU, 0x67d9bebeU, 0x724b3939U, + 0x94de4a4aU, 0x98d44c4cU, 0xb0e85858U, 0x854acfcfU, + 0xbb6bd0d0U, 0xc52aefefU, 0x4fe5aaaaU, 0xed16fbfbU, + 0x86c54343U, 0x9ad74d4dU, 0x66553333U, 0x11948585U, + 0x8acf4545U, 0xe910f9f9U, 0x04060202U, 0xfe817f7fU, + 0xa0f05050U, 0x78443c3cU, 0x25ba9f9fU, 0x4be3a8a8U, + 0xa2f35151U, 0x5dfea3a3U, 0x80c04040U, 0x058a8f8fU, + 0x3fad9292U, 0x21bc9d9dU, 0x70483838U, 0xf104f5f5U, + 0x63dfbcbcU, 0x77c1b6b6U, 0xaf75dadaU, 0x42632121U, + 0x20301010U, 0xe51affffU, 0xfd0ef3f3U, 0xbf6dd2d2U, + 0x814ccdcdU, 0x18140c0cU, 0x26351313U, 0xc32fececU, + 0xbee15f5fU, 0x35a29797U, 0x88cc4444U, 0x2e391717U, + 0x9357c4c4U, 0x55f2a7a7U, 0xfc827e7eU, 0x7a473d3dU, + 0xc8ac6464U, 0xbae75d5dU, 0x322b1919U, 0xe6957373U, + 0xc0a06060U, 0x19988181U, 0x9ed14f4fU, 0xa37fdcdcU, + 0x44662222U, 0x547e2a2aU, 0x3bab9090U, 0x0b838888U, + 0x8cca4646U, 0xc729eeeeU, 0x6bd3b8b8U, 0x283c1414U, + 0xa779dedeU, 0xbce25e5eU, 0x161d0b0bU, 0xad76dbdbU, + 0xdb3be0e0U, 0x64563232U, 0x744e3a3aU, 0x141e0a0aU, + 0x92db4949U, 0x0c0a0606U, 0x486c2424U, 0xb8e45c5cU, + 0x9f5dc2c2U, 0xbd6ed3d3U, 0x43efacacU, 0xc4a66262U, + 0x39a89191U, 0x31a49595U, 0xd337e4e4U, 0xf28b7979U, + 0xd532e7e7U, 0x8b43c8c8U, 0x6e593737U, 0xdab76d6dU, + 0x018c8d8dU, 0xb164d5d5U, 0x9cd24e4eU, 0x49e0a9a9U, + 0xd8b46c6cU, 0xacfa5656U, 0xf307f4f4U, 0xcf25eaeaU, + 0xcaaf6565U, 0xf48e7a7aU, 0x47e9aeaeU, 0x10180808U, + 0x6fd5babaU, 0xf0887878U, 0x4a6f2525U, 0x5c722e2eU, + 0x38241c1cU, 0x57f1a6a6U, 0x73c7b4b4U, 0x9751c6c6U, + 0xcb23e8e8U, 0xa17cddddU, 0xe89c7474U, 0x3e211f1fU, + 0x96dd4b4bU, 0x61dcbdbdU, 0x0d868b8bU, 0x0f858a8aU, + 0xe0907070U, 0x7c423e3eU, 0x71c4b5b5U, 0xccaa6666U, + 0x90d84848U, 0x06050303U, 0xf701f6f6U, 0x1c120e0eU, + 0xc2a36161U, 0x6a5f3535U, 0xaef95757U, 0x69d0b9b9U, + 0x17918686U, 0x9958c1c1U, 0x3a271d1dU, 0x27b99e9eU, + 0xd938e1e1U, 0xeb13f8f8U, 0x2bb39898U, 0x22331111U, + 0xd2bb6969U, 0xa970d9d9U, 0x07898e8eU, 0x33a79494U, + 0x2db69b9bU, 0x3c221e1eU, 0x15928787U, 0xc920e9e9U, + 0x8749ceceU, 0xaaff5555U, 0x50782828U, 0xa57adfdfU, + 0x038f8c8cU, 0x59f8a1a1U, 0x09808989U, 0x1a170d0dU, + 0x65dabfbfU, 0xd731e6e6U, 0x84c64242U, 0xd0b86868U, + 0x82c34141U, 0x29b09999U, 0x5a772d2dU, 0x1e110f0fU, + 0x7bcbb0b0U, 0xa8fc5454U, 0x6dd6bbbbU, 0x2c3a1616U +}; + +#define t_fn0(x) (sharedMemory[ (x)]) +#define t_fn1(x) (sharedMemory[256 + (x)]) +#define t_fn2(x) (sharedMemory[512 + (x)]) +#define t_fn3(x) (sharedMemory[768 + (x)]) + +__device__ __forceinline__ static void cn_aes_gpu_init(uint32_t *sharedMemory) +{ + for(int i = threadIdx.x; i < 1024; i += blockDim.x) + sharedMemory[i] = d_t_fn[i]; +} + +)===" +R"===( + +template< typename T > +__forceinline__ __device__ void unusedVar( const T& ) +{ +} + +template +__forceinline__ __device__ uint32_t shuffle(volatile uint32_t* ptr,const uint32_t sub,const int val,const uint32_t src) +{ +# if ( __CUDA_ARCH__ < 300 ) + ptr[sub] = val; + return ptr[src & (group_n-1)]; +# else + unusedVar( ptr ); + unusedVar( sub ); +# if (__CUDACC_VER_MAJOR__ >= 9) + return __shfl_sync(__activemask(), val, src, group_n); +# else + return __shfl( val, src, group_n ); +# endif +# endif +} + + +template +__forceinline__ __device__ uint64_t shuffle64(volatile uint32_t* ptr,const uint32_t sub,const uint64_t val,const uint32_t src, const uint32_t src2) +{ + uint64_t tmp; + ((uint32_t*)&tmp)[0] = shuffle(ptr, sub, static_cast(val), src); + ((uint32_t*)&tmp)[1] = shuffle(ptr, sub, static_cast(val >> 32), src2); + return tmp; +} + +struct u64 : public uint2 +{ + + __forceinline__ __device__ u64(){} + + __forceinline__ __device__ u64( const uint32_t x0, const uint32_t x1) + { + uint2::x = x0; + uint2::y = x1; + } + + __forceinline__ __device__ operator uint64_t() const + { + return *((uint64_t*)this); + } + + __forceinline__ __device__ u64( const uint64_t x0) + { + ((uint64_t*)&this->x)[0] = x0; + } + + __forceinline__ __device__ u64 operator^=(const u64& other) + { + uint2::x ^= other.x; + uint2::y ^= other.y; + + return *this; + } + + __forceinline__ __device__ u64 operator^=(const uint64_t& other) + { + uint2::x ^= static_cast(other); + uint2::y ^= static_cast(other >> 32); + + return *this; + } + + __forceinline__ __device__ u64 operator+(const u64& other) const + { + u64 tmp; + ((uint64_t*)&tmp.x)[0] = ((uint64_t*)&(this->x))[0] + ((uint64_t*)&(other.x))[0]; + + return tmp; + } + + __forceinline__ __device__ u64 operator+=(const uint64_t& other) + { + return ((uint64_t*)&this->x)[0] += other; + } +}; + +#ifdef RANDOM_MATH_64_BIT + +__device__ __forceinline__ static uint64_t rotate_left(uint64_t a, uint64_t b) +{ + const int shift = b & 63; + return (a << shift) | (a >> (64 - shift)); +} + +__device__ __forceinline__ static uint64_t rotate_right(uint64_t a, uint64_t b) +{ + const int shift = b & 63; + return (a >> shift) | (a << (64 - shift)); +} + +#else + +__device__ __forceinline__ static uint32_t rotate_left(uint32_t a, uint32_t b) { +#if __CUDA_ARCH__ < 350 + const uint32_t shift = b & 31; + return (a << shift) | (a >> (32 - shift)); +#else + return __funnelshift_l(a, a, b); +#endif +} +__device__ __forceinline__ static uint32_t rotate_right(uint32_t a, uint32_t b) { +#if __CUDA_ARCH__ < 350 + const uint32_t shift = b & 31; + return (a >> shift) | (a << (32 - shift)); +#else + return __funnelshift_r(a, a, b); +#endif +} + +#endif + +__global__ void CryptonightR_phase2( + const uint32_t ITERATIONS, + const size_t MEMORY, + const uint32_t MASK, + int threads, + int bfactor, + int partidx, + uint32_t *d_long_state, + uint32_t *d_ctx_a, + uint32_t *d_ctx_b, + uint32_t * d_ctx_state, + uint32_t startNonce, + uint32_t * __restrict__ d_input + ) +{ + __shared__ uint32_t sharedMemory[1024]; + + cn_aes_gpu_init( sharedMemory ); + +# if( __CUDA_ARCH__ < 300 ) + extern __shared__ uint64_t externShared[]; + // 8 x 64bit values + volatile uint64_t* myChunks = (volatile uint64_t*)(externShared + (threadIdx.x >> 1) * 8); + volatile uint32_t* sPtr = (volatile uint32_t*)(externShared + (blockDim.x >> 1) * 8) + (threadIdx.x & 0xFFFFFFFE); +# else + extern __shared__ uint64_t chunkMem[]; + volatile uint32_t* sPtr = NULL; + // 8 x 64bit values + volatile uint64_t* myChunks = (volatile uint64_t*)(chunkMem + (threadIdx.x >> 1) * 8); +# endif + + __syncthreads( ); + + const uint64_t tid = (blockDim.x * blockIdx.x + threadIdx.x); + const uint32_t thread = tid >> 1; + const uint32_t sub = tid & 1; + + if (thread >= threads) { + return; + } + + uint8_t *l0 = (uint8_t*)&d_long_state[((uint64_t)thread) * MEMORY]; + uint64_t ax0 = ((uint64_t*)(d_ctx_a + thread * 4))[sub]; + uint32_t idx0 = shuffle<2>(sPtr, sub, static_cast(ax0), 0); + uint64_t bx0 = ((uint64_t*)(d_ctx_b + thread * 16))[sub]; + uint64_t bx1 = ((uint64_t*)(d_ctx_b + thread * 16 + 4))[sub]; + + uint32_t r0 = d_ctx_b[thread * 16 + 4 * 2]; + uint32_t r1 = d_ctx_b[thread * 16 + 4 * 2 + 1]; + uint32_t r2 = d_ctx_b[thread * 16 + 4 * 2 + 2]; + uint32_t r3 = d_ctx_b[thread * 16 + 4 * 2 + 3]; + + const int batchsize = (ITERATIONS * 2) >> ( 1 + bfactor ); + const int start = partidx * batchsize; + const int end = start + batchsize; + + uint64_t* ptr0; + for (int i = start; i < end; ++i) { + ptr0 = (uint64_t *)&l0[idx0 & MASK & 0x1FFFC0]; + + ((ulonglong4*)myChunks)[sub] = ((ulonglong4*)ptr0)[sub]; + + uint32_t idx1 = (idx0 & 0x30) >> 3; + const u64 cx = myChunks[ idx1 + sub ]; + const u64 cx2 = myChunks[ idx1 + ((sub + 1) & 1) ]; + + u64 cx_aes = ax0 ^ u64( + t_fn0( cx.x & 0xff ) ^ t_fn1( (cx.y >> 8) & 0xff ) ^ t_fn2( (cx2.x >> 16) & 0xff ) ^ t_fn3( (cx2.y >> 24 ) ), + t_fn0( cx.y & 0xff ) ^ t_fn1( (cx2.x >> 8) & 0xff ) ^ t_fn2( (cx2.y >> 16) & 0xff ) ^ t_fn3( (cx.x >> 24 ) ) + ); + + { + const uint64_t chunk1 = myChunks[idx1 ^ 2 + sub]; + const uint64_t chunk2 = myChunks[idx1 ^ 4 + sub]; + const uint64_t chunk3 = myChunks[idx1 ^ 6 + sub]; + +#if(ALGO == cryptonight_r) + cx_aes ^= chunk1 ^ chunk2 ^ chunk3; +#endif + +#if (__CUDACC_VER_MAJOR__ >= 9) + __syncwarp(); +#else + __syncthreads(); +#endif + + myChunks[idx1 ^ 2 + sub] = chunk3 + bx1; + myChunks[idx1 ^ 4 + sub] = chunk1 + bx0; + myChunks[idx1 ^ 6 + sub] = chunk2 + ax0; + } + + myChunks[idx1 + sub] = cx_aes ^ bx0; + + ((ulonglong4*)ptr0)[sub] = ((ulonglong4*)myChunks)[sub]; + + idx0 = shuffle<2>(sPtr, sub, cx_aes.x, 0); + idx1 = (idx0 & 0x30) >> 3; + ptr0 = (uint64_t *)&l0[idx0 & MASK & 0x1FFFC0]; + + ((ulonglong4*)myChunks)[sub] = ((ulonglong4*)ptr0)[sub]; + + uint64_t cx_mul; + ((uint32_t*)&cx_mul)[0] = shuffle<2>(sPtr, sub, cx_aes.x , 0); + ((uint32_t*)&cx_mul)[1] = shuffle<2>(sPtr, sub, cx_aes.y , 0); + + const uint32_t r4 = shuffle<2>(sPtr, sub, static_cast(ax0), 0); + const uint32_t r6 = shuffle<2>(sPtr, sub, static_cast(bx0), 0); + const uint32_t r7 = shuffle<2>(sPtr, sub, static_cast(bx1), 0); + + const uint64_t ax0_saved = ax0; + + if (sub == 1) + { + ((uint32_t*)&myChunks[idx1])[0] ^= r0 + r1; + ((uint32_t*)&myChunks[idx1])[1] ^= r2 + r3; + + const uint32_t r5 = static_cast(ax0); +#if(ALGO == cryptonight_r) + const uint32_t r8 = static_cast(bx1); +#endif + + XMRSTAK_INCLUDE_RANDOM_MATH + } + +#if(ALGO == cryptonight_r) + r0 = shuffle<2>(sPtr, sub, r0, 1); + r1 = shuffle<2>(sPtr, sub, r1, 1); + r2 = shuffle<2>(sPtr, sub, r2, 1); + r3 = shuffle<2>(sPtr, sub, r3, 1); + ax0 ^= (sub == 0) ? (r2 | ((uint64_t)(r3) << 32)) : (r0 | ((uint64_t)(r1) << 32)); +#endif + +#if (__CUDACC_VER_MAJOR__ >= 9) + __syncwarp(); +#else + __syncthreads( ); +#endif + + uint64_t c = ((uint64_t*)myChunks)[idx1 + sub]; + + { + uint64_t cl = ((uint64_t*)myChunks)[idx1]; + // sub 0 -> hi, sub 1 -> lo + uint64_t res = sub == 0 ? __umul64hi( cx_mul, cl ) : cx_mul * cl; + + const uint64_t chunk1 = myChunks[ idx1 ^ 2 + sub ] +#if(ALGO == cryptonight_r_wow) + ^ res +#endif + ; + uint64_t chunk2 = myChunks[ idx1 ^ 4 + sub ]; +#if(ALGO == cryptonight_r_wow) + res ^= ((uint64_t*)&chunk2)[0]; +#endif + const uint64_t chunk3 = myChunks[ idx1 ^ 6 + sub ]; + +#if(ALGO == cryptonight_r) + cx_aes ^= chunk1 ^ chunk2 ^ chunk3; +#endif + +# if (__CUDACC_VER_MAJOR__ >= 9) + __syncwarp(); +# else + __syncthreads( ); +# endif + + myChunks[idx1 ^ 2 + sub] = chunk3 + bx1; + myChunks[idx1 ^ 4 + sub] = chunk1 + bx0; + myChunks[idx1 ^ 6 + sub] = chunk2 + ax0_saved; + + ax0 += res; + } + + bx1 = bx0; + bx0 = cx_aes; + + myChunks[idx1 + sub] = ax0; + + ((ulonglong4*)ptr0)[sub] = ((ulonglong4*)myChunks)[sub]; + + ax0 ^= c; + idx0 = shuffle<2>(sPtr, sub, static_cast(ax0), 0); + } + + if (bfactor > 0) + { + ((uint64_t*)(d_ctx_a + thread * 4))[sub] = ax0; + ((uint64_t*)(d_ctx_b + thread * 16))[sub] = bx0; + ((uint64_t*)(d_ctx_b + thread * 16 + 4))[sub] = bx1; + + if (sub == 1) + { + // must be valid only for `sub == 1` + d_ctx_b[thread * 16 + 4 * 2] = r0; + d_ctx_b[thread * 16 + 4 * 2 + 1] = r1; + d_ctx_b[thread * 16 + 4 * 2 + 2] = r2; + d_ctx_b[thread * 16 + 4 * 2 + 3] = r3; + } + } +} +)===" diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_device.hpp b/xmrstak/backend/nvidia/nvcc_code/cuda_device.hpp index 563bb3b9e..96cb679f5 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_device.hpp +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_device.hpp @@ -22,6 +22,18 @@ } \ ( (void) 0 ) +#define CU_CHECK(id, ...) { \ + CUresult result = __VA_ARGS__; \ + if(result != CUDA_SUCCESS){ \ + const char* s; \ + cuGetErrorString(result, &s); \ + std::cerr << "[CUDA] Error gpu " << id << ": <" << __FUNCTION__ << ">:" << __LINE__ << " \"" << (s ? s : "unknown error") << "\"" << std::endl; \ + throw std::runtime_error(std::string("[CUDA] Error: ") + std::string(s ? s : "unknown error")); \ + } \ +} \ +( (void) 0 ) + + /** execute and check a CUDA api command * * @param id gpu id (thread id) diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu index 43e21fb42..7a9ccddc2 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu @@ -129,14 +129,23 @@ __global__ void cryptonight_extra_gpu_prepare( int threads, uint32_t * __restric memcpy( d_ctx_a + thread * 4, ctx_a, 4 * 4 ); if(ALGO == cryptonight_monero_v8) { - memcpy( d_ctx_b + thread * 12, ctx_b, 4 * 4 ); + memcpy( d_ctx_b + thread * 16, ctx_b, 4 * 4 ); // bx1 XOR_BLOCKS_DST( ctx_state + 16, ctx_state + 20, ctx_b ); - memcpy( d_ctx_b + thread * 12 + 4, ctx_b, 4 * 4 ); + memcpy( d_ctx_b + thread * 16 + 4, ctx_b, 4 * 4 ); // division_result - memcpy( d_ctx_b + thread * 12 + 2 * 4, ctx_state + 24, 4 * 2 ); + memcpy( d_ctx_b + thread * 16 + 2 * 4, ctx_state + 24, 4 * 2 ); // sqrt_result - memcpy( d_ctx_b + thread * 12 + 2 * 4 + 2, ctx_state + 26, 4 * 2 ); + memcpy( d_ctx_b + thread * 16 + 2 * 4 + 2, ctx_state + 26, 4 * 2 ); + } + else if(ALGO == cryptonight_r_wow || ALGO == cryptonight_r) + { + memcpy(d_ctx_b + thread * 16, ctx_b, 4 * 4); + // bx1 + XOR_BLOCKS_DST(ctx_state + 16, ctx_state + 20, ctx_b); + memcpy(d_ctx_b + thread * 16 + 4, ctx_b, 4 * 4); + // r0, r1, r2, r3 + memcpy(d_ctx_b + thread * 16 + 2 * 4, ctx_state + 24, 4 * 8); } else memcpy( d_ctx_b + thread * 4, ctx_b, 4 * 4 ); @@ -258,6 +267,9 @@ extern "C" void cryptonight_extra_cpu_set_data( nvid_ctx* ctx, const void *data, extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx) { + CU_CHECK(ctx->device_id, cuDeviceGet(&ctx->cuDevice, ctx->device_id)); + CU_CHECK(ctx->device_id, cuCtxCreate(&ctx->cuContext, 0, ctx->cuDevice)); + cudaError_t err; err = cudaSetDevice(ctx->device_id); if(err != cudaSuccess) @@ -316,8 +328,16 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx) } else if(std::find(neededAlgorithms.begin(), neededAlgorithms.end(), cryptonight_monero_v8) != neededAlgorithms.end()) { - // bx1 (16byte), division_result (8byte) and sqrt_result (8byte) - ctx_b_size = 3 * 4 * sizeof(uint32_t) * wsize; + // bx0 (16byte), bx1 (16byte), division_result (8byte) and sqrt_result (8byte), padding (16byte) + ctx_b_size = 4 * 4 * sizeof(uint32_t) * wsize; + } + else if( + std::find(neededAlgorithms.begin(), neededAlgorithms.end(), cryptonight_r) != neededAlgorithms.end() || + std::find(neededAlgorithms.begin(), neededAlgorithms.end(), cryptonight_r_wow) != neededAlgorithms.end() + ) + { + // bx0 (16byte), bx1 (16byte), and [r0, r1, r2, r3] (a 8byte) + ctx_b_size = 4 * 4 * sizeof(uint32_t) * wsize; } else ctx->d_ctx_state2 = ctx->d_ctx_state; @@ -376,6 +396,16 @@ extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_extra_gpu_prepare<<>>( wsize, ctx->d_input, ctx->inputlen, startNonce, ctx->d_ctx_state,ctx->d_ctx_state2, ctx->d_ctx_a, ctx->d_ctx_b, ctx->d_ctx_key1, ctx->d_ctx_key2 )); } + else if(miner_algo == cryptonight_r) + { + CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_extra_gpu_prepare<<>>( wsize, ctx->d_input, ctx->inputlen, startNonce, + ctx->d_ctx_state,ctx->d_ctx_state2, ctx->d_ctx_a, ctx->d_ctx_b, ctx->d_ctx_key1, ctx->d_ctx_key2 )); + } + else if(miner_algo == cryptonight_r_wow) + { + CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_extra_gpu_prepare<<>>( wsize, ctx->d_input, ctx->inputlen, startNonce, + ctx->d_ctx_state,ctx->d_ctx_state2, ctx->d_ctx_a, ctx->d_ctx_b, ctx->d_ctx_key1, ctx->d_ctx_key2 )); + } else { /* pass two times d_ctx_state because the second state is used later in phase1, @@ -744,9 +774,11 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) // check if cryptonight_monero_v8 is selected for the user pool bool useCryptonight_v8 = (std::find(neededAlgorithms.begin(), neededAlgorithms.end(), cryptonight_monero_v8) != neededAlgorithms.end()); + bool useCryptonight_r = (std::find(neededAlgorithms.begin(), neededAlgorithms.end(), cryptonight_r) != neededAlgorithms.end()); + bool useCryptonight_r_wow = (std::find(neededAlgorithms.begin(), neededAlgorithms.end(), cryptonight_r_wow) != neededAlgorithms.end()); // overwrite default config if cryptonight_monero_v8 is mined and GPU has at least compute capability 5.0 - if(useCryptonight_v8 && gpuArch >= 50) + if((useCryptonight_v8 || useCryptonight_r || useCryptonight_r_wow) && gpuArch >= 50) { // 4 based on my test maybe it must be adjusted later size_t threads = 4; diff --git a/xmrstak/cli/cli-miner.cpp b/xmrstak/cli/cli-miner.cpp index 7539706e0..418726208 100644 --- a/xmrstak/cli/cli-miner.cpp +++ b/xmrstak/cli/cli-miner.cpp @@ -865,13 +865,12 @@ int do_benchmark(int block_version, int wait_sec, int work_sec) /* AMD and NVIDIA is currently only supporting work sizes up to 84byte * \todo fix this issue */ - xmrstak::miner_work benchWork = xmrstak::miner_work("", work, 84, 0, false, 0); printer::inst()->print_msg(L0, "Start a %d second benchmark...",work_sec); - xmrstak::globalStates::inst().switch_work(benchWork, dat); + xmrstak::globalStates::inst().switch_work(xmrstak::miner_work("", work, 84, 0, false, 0, 0), dat); uint64_t iStartStamp = get_timestamp_ms(); std::this_thread::sleep_for(std::chrono::seconds(work_sec)); - xmrstak::globalStates::inst().switch_work(oWork, dat); + xmrstak::globalStates::inst().switch_work(xmrstak::miner_work("", work, 84, 0, false, 0, 0), dat); double fTotalHps = 0.0; for (uint32_t i = 0; i < pvThreads->size(); i++) diff --git a/xmrstak/config.tpl b/xmrstak/config.tpl index 73ae054c2..c5fd882f0 100644 --- a/xmrstak/config.tpl +++ b/xmrstak/config.tpl @@ -25,11 +25,12 @@ R"===(// generated by XMRSTAK_VERSION * performance monitors, there is very little reason to spew out pages of text instead of concise reports. * Press 'h' (hashrate), 'r' (results) or 'c' (connection) to print reports. * - * verbose_level - 0 - Don't print anything. - * 1 - Print intro, connection event, disconnect event - * 2 - All of level 1, and new job (block) event if the difficulty is different from the last job - * 3 - All of level 1, and new job (block) event in all cases, result submission event. - * 4 - All of level 3, and automatic hashrate report printing + * verbose_level - 0 - Don't print anything. + * 1 - Print intro, connection event, disconnect event + * 2 - All of level 1, and new job (block) event if the difficulty is different from the last job + * 3 - All of level 1, and new job (block) event in all cases, result submission event. + * 4 - All of level 3, and automatic hashrate report printing + * 10 - Debug level for developer * * print_motd - Display messages from your pool operator in the hashrate result. */ diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp index 5dbddb09b..e60420234 100644 --- a/xmrstak/jconf.cpp +++ b/xmrstak/jconf.cpp @@ -87,38 +87,39 @@ constexpr size_t iConfigCnt = (sizeof(oConfigValues)/sizeof(oConfigValues[0])); xmrstak::coin_selection coins[] = { // name, userpool, devpool, default_pool_suggestion - { "aeon7", {POW(cryptonight_aeon)}, {POW(cryptonight_aeon)}, "mine.aeon-pool.com:5555" }, - { "bbscoin", {POW(cryptonight_aeon)}, {POW(cryptonight_aeon)}, nullptr }, - { "bittube", {POW(cryptonight_bittube2)}, {POW(cryptonight_heavy)}, "mining.bit.tube:13333" }, - { "cryptonight", {POW(cryptonight)}, {POW(cryptonight_monero_v8)}, nullptr }, - { "cryptonight_bittube2", {POW(cryptonight_bittube2)}, {POW(cryptonight_heavy)}, nullptr }, - { "cryptonight_masari", {POW(cryptonight_masari)}, {POW(cryptonight_monero_v8)}, nullptr }, - { "cryptonight_haven", {POW(cryptonight_haven)}, {POW(cryptonight_heavy)}, nullptr }, - { "cryptonight_heavy", {POW(cryptonight_heavy)}, {POW(cryptonight_heavy)}, nullptr }, + { "aeon7", {POW(cryptonight_aeon)}, {POW(cryptonight_aeon)}, "mine.aeon-pool.com:5555" }, + { "bbscoin", {POW(cryptonight_aeon)}, {POW(cryptonight_aeon)}, nullptr }, + { "bittube", {POW(cryptonight_bittube2)}, {POW(cryptonight_gpu)}, "mining.bit.tube:13333" }, + { "cryptonight", {POW(cryptonight)}, {POW(cryptonight_gpu)}, nullptr }, + { "cryptonight_bittube2", {POW(cryptonight_bittube2)}, {POW(cryptonight_gpu)}, nullptr }, + { "cryptonight_masari", {POW(cryptonight_masari)}, {POW(cryptonight_gpu)}, nullptr }, + { "cryptonight_haven", {POW(cryptonight_haven)}, {POW(cryptonight_gpu)}, nullptr }, + { "cryptonight_heavy", {POW(cryptonight_heavy)}, {POW(cryptonight_gpu)}, nullptr }, { "cryptonight_lite", {POW(cryptonight_lite)}, {POW(cryptonight_aeon)}, nullptr }, { "cryptonight_lite_v7", {POW(cryptonight_aeon)}, {POW(cryptonight_aeon)}, nullptr }, { "cryptonight_lite_v7_xor", {POW(cryptonight_ipbc)}, {POW(cryptonight_aeon)}, nullptr }, - { "cryptonight_superfast", {POW(cryptonight_superfast)}, {POW(cryptonight_monero_v8)}, nullptr }, + { "cryptonight_r", {POW(cryptonight_r)}, {POW(cryptonight_r),10,POW(cryptonight_monero_v8)}, nullptr }, + { "cryptonight_superfast", {POW(cryptonight_superfast)}, {POW(cryptonight_gpu)}, nullptr }, { "cryptonight_turtle", {POW(cryptonight_turtle)}, {POW(cryptonight_turtle)}, nullptr }, - { "cryptonight_v7", {POW(cryptonight_monero)}, {POW(cryptonight_monero_v8)}, nullptr }, - { "cryptonight_v8", {POW(cryptonight_monero_v8)}, {POW(cryptonight_monero_v8)}, nullptr }, - { "cryptonight_v8_half", {POW(cryptonight_v8_half)}, {POW(cryptonight_monero_v8)}, nullptr }, - { "cryptonight_v8_zelerius", {POW(cryptonight_v8_zelerius)},{POW(cryptonight_monero_v8)}, nullptr }, - { "cryptonight_v7_stellite", {POW(cryptonight_stellite)}, {POW(cryptonight_monero_v8)}, nullptr }, + { "cryptonight_v7", {POW(cryptonight_monero)}, {POW(cryptonight_gpu)}, nullptr }, + { "cryptonight_v8", {POW(cryptonight_monero_v8)}, {POW(cryptonight_r),10,POW(cryptonight_monero_v8)}, nullptr }, + { "cryptonight_v8_half", {POW(cryptonight_v8_half)}, {POW(cryptonight_gpu)}, nullptr }, + { "cryptonight_v8_zelerius", {POW(cryptonight_v8_zelerius)},{POW(cryptonight_gpu)}, nullptr }, + { "cryptonight_v7_stellite", {POW(cryptonight_stellite)}, {POW(cryptonight_gpu)}, nullptr }, { "cryptonight_gpu", {POW(cryptonight_gpu)}, {POW(cryptonight_gpu)}, "pool.ryo-currency.com:3333" }, - { "cryptonight_conceal", {POW(cryptonight_conceal)}, {POW(cryptonight_monero_v8)}, nullptr }, - { "freehaven", {POW(cryptonight_superfast)}, {POW(cryptonight_monero_v8)}, nullptr }, - { "graft", {POW(cryptonight_monero_v8)}, {POW(cryptonight_monero_v8)}, nullptr }, - { "haven", {POW(cryptonight_haven)}, {POW(cryptonight_heavy)}, nullptr }, - { "lethean", {POW(cryptonight_monero)}, {POW(cryptonight_monero_v8)}, nullptr }, - { "masari", {POW(cryptonight_v8_half)}, {POW(cryptonight_monero_v8)}, nullptr }, - { "monero", {POW(cryptonight_monero_v8)}, {POW(cryptonight_monero_v8)}, "pool.usxmrpool.com:3333" }, - { "qrl", {POW(cryptonight_monero)}, {POW(cryptonight_monero_v8)}, nullptr }, - { "ryo", {POW(cryptonight_gpu)}, {POW(cryptonight_gpu)}, "pool.ryo-currency.com:3333" }, - { "stellite", {POW(cryptonight_v8_half)}, {POW(cryptonight_monero_v8)}, nullptr }, + { "cryptonight_conceal", {POW(cryptonight_conceal)}, {POW(cryptonight_gpu)}, nullptr }, + { "freehaven", {POW(cryptonight_superfast)}, {POW(cryptonight_gpu)}, nullptr }, + { "graft", {POW(cryptonight_monero_v8)}, {POW(cryptonight_gpu)}, nullptr }, + { "haven", {POW(cryptonight_haven)}, {POW(cryptonight_gpu)}, nullptr }, + { "lethean", {POW(cryptonight_monero)}, {POW(cryptonight_gpu)}, nullptr }, + { "masari", {POW(cryptonight_v8_half)}, {POW(cryptonight_gpu)}, nullptr }, + { "monero", {POW(cryptonight_r),10,POW(cryptonight_monero_v8)}, {POW(cryptonight_r),10,POW(cryptonight_monero_v8)}, "pool.usxmrpool.com:3333" }, + { "qrl", {POW(cryptonight_monero)}, {POW(cryptonight_gpu)}, nullptr }, + { "ryo", {POW(cryptonight_gpu)}, {POW(cryptonight_gpu)}, "pool.ryo-currency.com:3333" }, + { "stellite", {POW(cryptonight_v8_half)}, {POW(cryptonight_gpu)}, nullptr }, { "turtlecoin", {POW(cryptonight_turtle), 6u,POW(cryptonight_aeon)}, {POW(cryptonight_aeon)}, nullptr }, { "plenteum", {POW(cryptonight_turtle)}, {POW(cryptonight_turtle)}, nullptr }, - { "zelerius", {POW(cryptonight_v8_zelerius), 7, POW(cryptonight_monero_v8)}, {POW(cryptonight_monero_v8)}, nullptr } + { "zelerius", {POW(cryptonight_v8_zelerius), 7, POW(cryptonight_monero_v8)}, {POW(cryptonight_gpu)}, nullptr } }; constexpr size_t coin_algo_size = (sizeof(coins)/sizeof(coins[0])); diff --git a/xmrstak/misc/console.hpp b/xmrstak/misc/console.hpp index 5d78772c3..6df6597c6 100644 --- a/xmrstak/misc/console.hpp +++ b/xmrstak/misc/console.hpp @@ -21,7 +21,7 @@ inline long long unsigned int int_port(size_t i) return i; } -enum verbosity : size_t { L0 = 0, L1 = 1, L2 = 2, L3 = 3, L4 = 4, LINF = 100}; +enum verbosity : size_t { L0 = 0, L1 = 1, L2 = 2, L3 = 3, L4 = 4, LDEBUG = 10, LINF = 100}; class printer { diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp index 0dd7db137..24e348897 100644 --- a/xmrstak/misc/executor.cpp +++ b/xmrstak/misc/executor.cpp @@ -125,9 +125,8 @@ bool executor::get_live_pools(std::vector& eval_pools, bool is_dev) if(xmrstak::globalStates::inst().pool_id != invalid_pool_id) { printer::inst()->print_msg(L0, "All pools are dead. Idling..."); - auto work = xmrstak::miner_work(); xmrstak::pool_data dat; - xmrstak::globalStates::inst().switch_work(work, dat); + xmrstak::globalStates::inst().switch_work(xmrstak::miner_work(), dat); } if(over_limit == pool_count) @@ -365,13 +364,12 @@ void executor::on_pool_have_job(size_t pool_id, pool_job& oPoolJob) jpsock* pool = pick_pool_by_id(pool_id); - xmrstak::miner_work oWork(oPoolJob.sJobID, oPoolJob.bWorkBlob, oPoolJob.iWorkLen, oPoolJob.iTarget, pool->is_nicehash(), pool_id); - xmrstak::pool_data dat; dat.iSavedNonce = oPoolJob.iSavedNonce; dat.pool_id = pool_id; - xmrstak::globalStates::inst().switch_work(oWork, dat); + xmrstak::globalStates::inst().switch_work(xmrstak::miner_work(oPoolJob.sJobID, oPoolJob.bWorkBlob, + oPoolJob.iWorkLen, oPoolJob.iTarget, pool->is_nicehash(), pool_id, oPoolJob.iBlockHeight), dat); if(dat.pool_id != pool_id) { @@ -446,7 +444,7 @@ void executor::on_miner_result(size_t pool_id, job_result& oResult) if(bResult) { uint64_t* targets = (uint64_t*)oResult.bResult; - log_result_ok(jpsock::t64_to_diff(targets[3])); + log_result_ok(t64_to_diff(targets[3])); printer::inst()->print_msg(L3, "Result accepted by the pool."); } else @@ -578,8 +576,13 @@ void executor::ex_main() else pools.emplace_front(0, "donate.xmr-stak.net:4444", "", "", "", 0.0, true, false, "", true); break; + case cryptonight_r: + if(dev_tls) + pools.emplace_front(0, "donate.xmr-stak.net:8822", "", "", "", 0.0, true, true, "", false); + else + pools.emplace_front(0, "donate.xmr-stak.net:5522", "", "", "", 0.0, true, false, "", false); + break; default: - case cryptonight_lite: if(dev_tls) pools.emplace_front(0, "donate.xmr-stak.net:6666", "", "", "", 0.0, true, true, "", false); else diff --git a/xmrstak/misc/jext.hpp b/xmrstak/misc/jext.hpp index f4a333c22..9936fa813 100644 --- a/xmrstak/misc/jext.hpp +++ b/xmrstak/misc/jext.hpp @@ -14,3 +14,49 @@ inline const Value* GetObjectMember(const Value& obj, const char* key) else return nullptr; } + +#ifdef _MSC_VER + +#include +#define bswap_32(x) _byteswap_ulong(x) +#define bswap_64(x) _byteswap_uint64(x) + +#elif defined(__APPLE__) + +// Mac OS X / Darwin features +#include +#define bswap_32(x) OSSwapInt32(x) +#define bswap_64(x) OSSwapInt64(x) + +#elif defined(__sun) || defined(sun) + +#include +#define bswap_32(x) BSWAP_32(x) +#define bswap_64(x) BSWAP_64(x) + +#elif defined(__FreeBSD__) + +#include +#define bswap_32(x) bswap32(x) +#define bswap_64(x) bswap64(x) + +#elif defined(__OpenBSD__) + +#include +#define bswap_32(x) swap32(x) +#define bswap_64(x) swap64(x) + +#elif defined(__NetBSD__) + +#include +#include +#if defined(__BSWAP_RENAME) && !defined(__bswap_32) +#define bswap_32(x) bswap32(x) +#define bswap_64(x) bswap64(x) +#endif + +#else + +#include + +#endif diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp index d5b0d7fcc..786b18b4f 100644 --- a/xmrstak/net/jpsock.cpp +++ b/xmrstak/net/jpsock.cpp @@ -403,11 +403,12 @@ bool jpsock::process_pool_job(const opq_json_val* params, const uint64_t message if (!params->val->IsObject()) return set_socket_error("PARSE error: Job error 1"); - const Value *blob, *jobid, *target, *motd; + const Value *blob, *jobid, *target, *motd, *blk_height; jobid = GetObjectMember(*params->val, "job_id"); blob = GetObjectMember(*params->val, "blob"); target = GetObjectMember(*params->val, "target"); motd = GetObjectMember(*params->val, "motd"); + blk_height = GetObjectMember(*params->val, "height"); if (jobid == nullptr || blob == nullptr || target == nullptr || !jobid->IsString() || !blob->IsString() || !target->IsString()) @@ -445,10 +446,8 @@ bool jpsock::process_pool_job(const opq_json_val* params, const uint64_t message // lock reading of oCurrentJob std::unique_lock jobIdLock(job_mutex); // compare possible non equal length job id's - if(iWorkLen == oCurrentJob.iWorkLen && - memcmp(oPoolJob.bWorkBlob, oCurrentJob.bWorkBlob, iWorkLen) == 0 && - strcmp(jobid->GetString(), oCurrentJob.sJobID) == 0 - ) + if(iWorkLen == oCurrentJob.iWorkLen && memcmp(oPoolJob.bWorkBlob, oCurrentJob.bWorkBlob, iWorkLen) == 0 && + strcmp(jobid->GetString(), oCurrentJob.sJobID) == 0) { return set_socket_error("Duplicate equal job detected! Please contact your pool admin."); } @@ -466,7 +465,6 @@ bool jpsock::process_pool_job(const opq_json_val* params, const uint64_t message if(!hex2bin(sTempStr, 8, (unsigned char*)&iTempInt) || iTempInt == 0) return set_socket_error("PARSE error: Invalid target"); - oPoolJob.iTarget = t32_to_t64(iTempInt); } else if(target_slen <= 16) @@ -481,6 +479,9 @@ bool jpsock::process_pool_job(const opq_json_val* params, const uint64_t message return set_socket_error("PARSE error: Job error 5"); iJobDiff = t64_to_diff(oPoolJob.iTarget); + + if(blk_height != nullptr && blk_height->IsUint64()) + oPoolJob.iBlockHeight = bswap_64(blk_height->GetUint64()); std::unique_lock lck(job_mutex); oCurrentJob = oPoolJob; diff --git a/xmrstak/net/jpsock.hpp b/xmrstak/net/jpsock.hpp index a1112df74..949764813 100644 --- a/xmrstak/net/jpsock.hpp +++ b/xmrstak/net/jpsock.hpp @@ -66,11 +66,6 @@ class jpsock std::string&& get_call_error(); bool have_call_error() { return call_error; } bool have_sock_error() { return bHaveSocketError; } - - inline static uint64_t t32_to_t64(uint32_t t) { return 0xFFFFFFFFFFFFFFFFULL / (0xFFFFFFFFULL / ((uint64_t)t)); } - inline static uint64_t t64_to_diff(uint64_t t) { return 0xFFFFFFFFFFFFFFFFULL / t; } - inline static uint64_t diff_to_t64(uint64_t d) { return 0xFFFFFFFFFFFFFFFFULL / d; } - inline uint64_t get_current_diff() { return iJobDiff; } void save_nonce(uint32_t nonce); diff --git a/xmrstak/net/msgstruct.hpp b/xmrstak/net/msgstruct.hpp index cd23a94c4..813fc7d06 100644 --- a/xmrstak/net/msgstruct.hpp +++ b/xmrstak/net/msgstruct.hpp @@ -16,6 +16,7 @@ struct pool_job uint64_t iTarget; uint32_t iWorkLen; uint32_t iSavedNonce; + uint64_t iBlockHeight = uint64_t(-1); pool_job() : iWorkLen(0), iSavedNonce(0) {} pool_job(const char* sJobID, uint64_t iTarget, const uint8_t* bWorkBlob, uint32_t iWorkLen) : @@ -175,6 +176,10 @@ struct ex_event } }; +inline uint64_t t32_to_t64(uint32_t t) { return 0xFFFFFFFFFFFFFFFFULL / (0xFFFFFFFFULL / ((uint64_t)t)); } +inline uint64_t t64_to_diff(uint64_t t) { return 0xFFFFFFFFFFFFFFFFULL / t; } +inline uint64_t diff_to_t64(uint64_t d) { return 0xFFFFFFFFFFFFFFFFULL / d; } + #include //Get steady_clock timestamp - misc helper function inline size_t get_timestamp()