diff --git a/.appveyor.yml b/.appveyor.yml
index c336842d5d34d5e2f3cc174b83bd4c853c029693..fee8f6524d41d2180f80663f2a61222e053b71d8 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 09ff7aef519146966af9b4a56e19930587e82c52..a5c06df8a20a735a9e6e1b7245953b2caa77ca95 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 0000000000000000000000000000000000000000..4aabe51d0bc4dbd6acea64b5eac9240ab278cd78
--- /dev/null
+++ b/xmrstak/backend/amd/OclCryptonightR_gen.cpp
@@ -0,0 +1,354 @@
+#include <string>
+#include <sstream>
+#include <mutex>
+#include <cstring>
+#include <thread>
+
+
+#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 <chrono>
+#include <thread>
+#include <iostream>
+
+
+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<typename T>
+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<CacheEntry> CryptonightR_cache;
+
+static std::mutex background_tasks_mutex;
+static std::vector<BackgroundTaskBase*> background_tasks;
+static std::thread* background_thread = nullptr;
+
+static void background_thread_proc()
+{
+ std::vector<BackgroundTaskBase*> tasks;
+ for (;;) {
+ tasks.clear();
+ {
+ std::lock_guard<std::mutex> 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<typename T>
+static void background_exec(T&& func)
+{
+ BackgroundTaskBase* task = new BackgroundTask<T>(std::move(func));
+
+ std::lock_guard<std::mutex> 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<cl_program> 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<std::mutex> 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<<BuildLog<<std::endl;
+
+ free(BuildLog);
+ return program;
+ }
+
+ cl_build_status status;
+ do
+ {
+ if((ret = clGetProgramBuildInfo(program, ctx->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<cryptonight_r_wow>(code, height);
+ break;
+ case cryptonight_r:
+ code_size = v4_random_math_init<cryptonight_r>(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 0000000000000000000000000000000000000000..a69df9074be2715035948358d9f2c0cdfe61515e
--- /dev/null
+++ b/xmrstak/backend/amd/OclCryptonightR_gen.hpp
@@ -0,0 +1,26 @@
+#pragma once
+
+#include "xmrstak/backend/cryptonight.hpp"
+
+#include <stdint.h>
+#include <vector>
+#include <string>
+
+#if defined(__APPLE__)
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#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 8713784c274e56a25213304c32959c1fc5d94f1c..a2cbe8f5467ad83d668d073bc6321c182e7333f9 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 <stdio.h>
#include <string.h>
@@ -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<std::shared_ptr<InterleaveData>> interleaveData(num_gpus, nullptr);
+ std::vector<cl_context> 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<double>(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 d665dff1b2f3d5a5296178fb37467caf64001c59..ae2b506dbeeb673ac68b21c2db8adb20cfde7b61 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<xmrstak_algo_id, cl_program> Program;
std::map<xmrstak_algo_id, std::array<cl_kernel,8>> 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<GpuContext> 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 0000000000000000000000000000000000000000..9edb774adcbc2a31bb25f3d6e581960d6b71a00f
--- /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 <http://www.gnu.org/licenses/>.
+ *
+ */
+
+#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 48f4ca49aa53019507d6ea7350f911ae2740e981..ea688e0534239547c7da41517201e06695851724 100644
--- a/xmrstak/backend/amd/autoAdjust.hpp
+++ b/xmrstak/backend/amd/autoAdjust.hpp
@@ -139,6 +139,10 @@ private:
// 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 @@ private:
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 1c9eb62797bf491e84cb3ca38a5cd82428fa7e7a..eb00094137ba75fde007da4da671b73b30d6fe85 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 5c9a73332fc1f0435e6eec210ffa3d73981fcc4a..a7c77cdac0d5dbd6b8164db4e753b655b1da049d 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 <stddef.h>
#include <inttypes.h>
-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 dc378e88a56d537e4db998f55b24b061e3965277..43f71987349f7f92b085012bf01ed83ee7ea1b20 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 <memory.h>
#include <stdio.h>
@@ -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<uint32_t>(al); \
+ cn_r_data[5] = static_cast<uint32_t>(ah); \
+ cn_r_data[6] = static_cast<uint32_t>(_mm_cvtsi128_si32(bx0)); \
+ cn_r_data[7] = static_cast<uint32_t>(_mm_cvtsi128_si32(bx1)); \
+ cn_r_data[8] = static_cast<uint32_t>(_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<N> 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<xmrstak_algo_id ALGO>
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<false, false, ALGO>((__m128i*)ctx[0]->hash_state, (__m128i*)ctx[0]->long_state, algo);
@@ -1087,7 +1124,6 @@ struct Cryptonight_hash_gpu
template<xmrstak_algo_id ALGO, bool SOFT_AES, bool PREFETCH>
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<PREFETCH, ALGO>(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<size_t N>
+struct Cryptonight_R_generator
+{
+ template<xmrstak_algo_id ALGO>
+ 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<ALGO>(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 0000000000000000000000000000000000000000..07dd3cf6192d5fbae4ee49b387fda3e76c207797
--- /dev/null
+++ b/xmrstak/backend/cpu/crypto/variant4_random_math.h
@@ -0,0 +1,451 @@
+#pragma once
+
+#include <string.h>
+#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<typename v4_reg>
+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<xmrstak_algo_id ALGO>
+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 50507f2aee1e0e3bcce67eae76cd7c1f8f8f293b..064b07339d4e1f7a6bb040d971cc1fc2787e6410 100644
--- a/xmrstak/backend/cpu/minethd.cpp
+++ b/xmrstak/backend/cpu/minethd.cpp
@@ -50,6 +50,7 @@
#include <cstring>
#include <thread>
#include <bitset>
+#include <unordered_map>
#ifdef _WIN32
#include <windows.h>
@@ -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<size_t N>
-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<N>::template hash<cryptonight_conceal, false, false>,
Cryptonight_hash<N>::template hash<cryptonight_conceal, true, false>,
Cryptonight_hash<N>::template hash<cryptonight_conceal, false, true>,
- Cryptonight_hash<N>::template hash<cryptonight_conceal, true, true>
+ Cryptonight_hash<N>::template hash<cryptonight_conceal, true, true>,
+
+ Cryptonight_hash<N>::template hash<cryptonight_r, false, false>,
+ Cryptonight_hash<N>::template hash<cryptonight_r, true, false>,
+ Cryptonight_hash<N>::template hash<cryptonight_r, false, true>,
+ Cryptonight_hash<N>::template hash<cryptonight_r, true, true>
};
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<cryptonight_monero_v8>;
+ hash_fun = Cryptonight_hash_asm<1u, 0u>::template hash<cryptonight_monero_v8>;
else if(N == 2)
- selected_function = Cryptonight_hash_asm<2u, 0u>::template hash<cryptonight_monero_v8>;
+ hash_fun = Cryptonight_hash_asm<2u, 0u>::template hash<cryptonight_monero_v8>;
}
// 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<cryptonight_monero_v8>;
+ hash_fun = Cryptonight_hash_asm<1u, 1u>::template hash<cryptonight_monero_v8>;
}
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<uint32_t, minethd::cn_on_new_job> on_new_job_map = {
+ {cryptonight_r, Cryptonight_R_generator<N>::template cn_on_new_job<cryptonight_r>},
+ };
+
+ 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<N>(::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<N>(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<N>(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo, asm_version_str);
+ func_multi_selector<N>(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<N>(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo, asm_version_str);
+ func_multi_selector<N>(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 41315269bac2390a11b5d9a3c240885a39e0c9f1..ca89e5b524d3589e36c221a9e960164878e7fa19 100644
--- a/xmrstak/backend/cpu/minethd.hpp
+++ b/xmrstak/backend/cpu/minethd.hpp
@@ -22,6 +22,7 @@ public:
static std::vector<iBackend*> 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 @@ public:
static cryptonight_ctx* minethd_alloc_ctx();
-private:
-
template<size_t N>
- 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<uint32_t N>
diff --git a/xmrstak/backend/cryptonight.hpp b/xmrstak/backend/cryptonight.hpp
index db07877897e8e6bd8ff78eb770056351fc9caa98..00311bb9333efa5964f5db487516dd088eebcbe8 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<std::string, 15> base_algo_names =
+ static std::array<std::string, 17> 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<std::string, 3> 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<xmrstak_algo, 15> pow = {{
+ static std::array<xmrstak_algo, 17> 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<xmrstak_algo, 3> derived_pow =
diff --git a/xmrstak/backend/globalStates.cpp b/xmrstak/backend/globalStates.cpp
index 4eeed3c4b832c74ecaba9a7266d363dc098e531c..52ef3f39150ea9bb9b4607c89fd867a15dfc764f 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 c8d691712ce0ad1eebf9321f798b8b5cca2b97f6..d6966c4a2ed68f91acd63f2c801aca2327eaabea 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 b6456f031186639b5278431b250663379bfe04df..c8174df328945e6241321e0f38cd4dcb5365e593 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 0000000000000000000000000000000000000000..87eb05540b83bde74ab295f8810b92194a03d131
--- /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 <http://www.gnu.org/licenses/>.
+ *
+ */
+
+#include <string>
+#include <sstream>
+#include <mutex>
+#include <cstring>
+#include <nvrtc.h>
+#include <thread>
+
+#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<char>& 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<char> ptx;
+ std::string lowered_name;
+};
+
+struct BackgroundTaskBase
+{
+ virtual ~BackgroundTaskBase() {}
+ virtual void exec() = 0;
+};
+
+template<typename T>
+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<CacheEntry> CryptonightR_cache;
+
+static std::mutex background_tasks_mutex;
+static std::vector<BackgroundTaskBase*> background_tasks;
+static std::thread* background_thread = nullptr;
+
+static void background_thread_proc()
+{
+ std::vector<BackgroundTaskBase*> tasks;
+ for (;;) {
+ tasks.clear();
+ {
+ std::lock_guard<std::mutex> 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<typename T>
+static void background_exec(T&& func)
+{
+ BackgroundTaskBase* task = new BackgroundTask<T>(std::move(func));
+
+ std::lock_guard<std::mutex> 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<char>& 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<std::mutex> 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<int>(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<char>& 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<char> 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<cryptonight_r_wow>(code, height);
+ break;
+ case cryptonight_r:
+ code_size = v4_random_math_init<cryptonight_r>(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 0000000000000000000000000000000000000000..e214647b97427c049aed29c3e4c6e6f66f0cdf66
--- /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 <http://www.gnu.org/licenses/>.
+ *
+ */
+
+#pragma once
+
+#include "xmrstak/backend/cryptonight.hpp"
+
+#include <stdint.h>
+#include <vector>
+#include <string>
+
+
+namespace xmrstak
+{
+namespace nvidia
+{
+
+void CryptonightR_get_program(std::vector<char>& 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 07ed4d31e9a118809bfd4361ab89eba09aee6907..794e68d112d0e23c85f7a52a81abb5eab7569381 100644
--- a/xmrstak/backend/nvidia/minethd.cpp
+++ b/xmrstak/backend/nvidia/minethd.cpp
@@ -173,6 +173,8 @@ std::vector<iBackend*>* 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 45ffef80668c5b1726025d5756dfa04312bcdb68..fe77b6f81486d85df2a05aa5174d16a0b968c684 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 <cuda.h>
+
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 184825222ce86cd94d57dd2b1747463586edbf77..d082f3362035d3acc4ede3c11d603bdedb36f20d 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<xmrstak_algo_id ALGO, uint32_t MEM_MODE>
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<ALGO, MEM_MODE><<<
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<char> 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_gpu, 1>,
cryptonight_core_gpu_hash<cryptonight_conceal, 0>,
- cryptonight_core_gpu_hash<cryptonight_conceal, 1>
+ cryptonight_core_gpu_hash<cryptonight_conceal, 1>,
+
+ cryptonight_core_gpu_hash<cryptonight_r_wow, 0>,
+ cryptonight_core_gpu_hash<cryptonight_r_wow, 1>,
+
+ cryptonight_core_gpu_hash<cryptonight_r, 0>,
+ cryptonight_core_gpu_hash<cryptonight_r, 1>
};
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 0000000000000000000000000000000000000000..bcf49508009f691fdafbdf48f14690c91ab501e2
--- /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 <http://www.gnu.org/licenses/>.
+ *
+ */
+
+#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<size_t group_n>
+__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<size_t group_n>
+__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<group_n>(ptr, sub, static_cast<uint32_t>(val), src);
+ ((uint32_t*)&tmp)[1] = shuffle<group_n>(ptr, sub, static_cast<uint32_t>(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<uint32_t>(other);
+ uint2::y ^= static_cast<uint32_t>(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<uint32_t>(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<uint32_t>(ax0), 0);
+ const uint32_t r6 = shuffle<2>(sPtr, sub, static_cast<uint32_t>(bx0), 0);
+ const uint32_t r7 = shuffle<2>(sPtr, sub, static_cast<uint32_t>(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<uint32_t>(ax0);
+#if(ALGO == cryptonight_r)
+ const uint32_t r8 = static_cast<uint32_t>(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<uint32_t>(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 563bb3b9e79c02e0718a8839589f87f9f90365a3..96cb679f5a165be4f3f82cea2a6d0ace88f47ee6 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 43e21fb428ce9d7772a740f0247dad596ce00b17..7a9ccddc222f296b1b653471a724f19dbf92070e 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<cryptonight_gpu><<<grid, block >>>( wsize, ctx->d_input, ctx->inputlen, startNonce,
ctx->d_ctx_state,ctx->d_ctx_state2, ctx->d_ctx_a, ctx->d_ctx_b, ctx->d_ctx_key1, ctx->d_ctx_key2 ));
}
+ else if(miner_algo == cryptonight_r)
+ {
+ CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_extra_gpu_prepare<cryptonight_r><<<grid, block >>>( wsize, ctx->d_input, ctx->inputlen, startNonce,
+ ctx->d_ctx_state,ctx->d_ctx_state2, ctx->d_ctx_a, ctx->d_ctx_b, ctx->d_ctx_key1, ctx->d_ctx_key2 ));
+ }
+ else if(miner_algo == cryptonight_r_wow)
+ {
+ CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_extra_gpu_prepare<cryptonight_r_wow><<<grid, block >>>( wsize, ctx->d_input, ctx->inputlen, startNonce,
+ ctx->d_ctx_state,ctx->d_ctx_state2, ctx->d_ctx_a, ctx->d_ctx_b, ctx->d_ctx_key1, ctx->d_ctx_key2 ));
+ }
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 7539706e0cb3619869a1357912b6bf471d38d0fa..418726208d8b273fc9fed85ad4ed6506131ed3c7 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 73ae054c2346661360022271b64e4ec9b2f48feb..d8fd861a7ccafe688473ffb58f9698ae97782d08 100644
--- a/xmrstak/config.tpl
+++ b/xmrstak/config.tpl
@@ -25,15 +25,16 @@ 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.
*/
-"verbose_level" : 3,
+"verbose_level" : 4,
"print_motd" : true,
/*
@@ -42,7 +43,7 @@ R"===(// generated by XMRSTAK_VERSION
* h_print_time - How often, in seconds, should we print a hashrate report if verbose_level is set to 4.
* This option has no effect if verbose_level is not 4.
*/
-"h_print_time" : 60,
+"h_print_time" : 300,
/*
* Manual hardware AES override
diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp
index 5dbddb09b4c3c3c47b8d4c7d10190c3cfe565c85..e60420234454eaa3f8c36a6db0be167031c69d00 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 5d78772c3ea87a094f41e9c180c799d3982fbb5e..6df6597c6e63fcb230d0f39555d486c59b0d612a 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 0dd7db1370c06a93f2aa365495fce38f626c5deb..24e3488977b78a94fe188830cab1ab0f6f89679d 100644
--- a/xmrstak/misc/executor.cpp
+++ b/xmrstak/misc/executor.cpp
@@ -125,9 +125,8 @@ bool executor::get_live_pools(std::vector<jpsock*>& 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 f4a333c220f23493666da340bf9473b8dfdcbc13..9936fa81329169a979f18c5cd99bf2a2f87c3382 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 <stdlib.h>
+#define bswap_32(x) _byteswap_ulong(x)
+#define bswap_64(x) _byteswap_uint64(x)
+
+#elif defined(__APPLE__)
+
+// Mac OS X / Darwin features
+#include <libkern/OSByteOrder.h>
+#define bswap_32(x) OSSwapInt32(x)
+#define bswap_64(x) OSSwapInt64(x)
+
+#elif defined(__sun) || defined(sun)
+
+#include <sys/byteorder.h>
+#define bswap_32(x) BSWAP_32(x)
+#define bswap_64(x) BSWAP_64(x)
+
+#elif defined(__FreeBSD__)
+
+#include <sys/endian.h>
+#define bswap_32(x) bswap32(x)
+#define bswap_64(x) bswap64(x)
+
+#elif defined(__OpenBSD__)
+
+#include <sys/types.h>
+#define bswap_32(x) swap32(x)
+#define bswap_64(x) swap64(x)
+
+#elif defined(__NetBSD__)
+
+#include <sys/types.h>
+#include <machine/bswap.h>
+#if defined(__BSWAP_RENAME) && !defined(__bswap_32)
+#define bswap_32(x) bswap32(x)
+#define bswap_64(x) bswap64(x)
+#endif
+
+#else
+
+#include <byteswap.h>
+
+#endif
diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp
index d5b0d7fcc6d18b5b9f968b0c5f870069c534dcbc..786b18b4f89c3b9d17e7ec5b8282dbe97c180414 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<std::mutex> 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<std::mutex> lck(job_mutex);
oCurrentJob = oPoolJob;
diff --git a/xmrstak/net/jpsock.hpp b/xmrstak/net/jpsock.hpp
index a1112df7430458f73294b2eff0537d7fbbb81a5f..94976481326fa0c8ba99daaebf48ae6d57fb70b9 100644
--- a/xmrstak/net/jpsock.hpp
+++ b/xmrstak/net/jpsock.hpp
@@ -66,11 +66,6 @@ public:
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 cd23a94c4d4ed2316bf31b0bc8c23c0447b29eca..813fc7d06ea6dfbf42fa99be0cca85a6a52aa2d5 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 <chrono>
//Get steady_clock timestamp - misc helper function
inline size_t get_timestamp()
diff --git a/xmrstak/version.cpp b/xmrstak/version.cpp
index a118989ccaa7aad10f3d8dee4444c04cbf57a50d..bcaebf4e0f9f193a667d816f4b1eee43b104bf93 100644
--- a/xmrstak/version.cpp
+++ b/xmrstak/version.cpp
@@ -18,7 +18,7 @@
#endif
#define XMR_STAK_NAME "xmr-stak"
-#define XMR_STAK_VERSION "2.8.3"
+#define XMR_STAK_VERSION "2.9.0"
#if defined(_WIN32)
#define OS_TYPE "win"