diff --git a/CMakeLists.txt b/CMakeLists.txt
index 15a2684ca18cec152038814768c56dbacb5add5e..784c0bde7e00e507bec0095a46f48ac92f82c57d 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -148,6 +148,11 @@ if(CUDA_ENABLE)
set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -std=c++11")
endif()
+ # avoid that nvcc in CUDA 8 complains about sm_20 pending removal
+ if(CUDA_VERSION VERSION_EQUAL 8.0)
+ set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -Wno-deprecated-gpu-targets")
+ endif()
+
# avoid that nvcc in CUDA < 8 tries to use libc `memcpy` within the kernel
if(CUDA_VERSION VERSION_LESS 8.0)
add_definitions(-D_FORCE_INLINES)
@@ -522,8 +527,8 @@ else()
add_executable(xmr-stak ${SRCFILES_CPP})
endif()
-set(EXECUTABLE_OUTPUT_PATH "bin")
-set(LIBRARY_OUTPUT_PATH "bin")
+set(EXECUTABLE_OUTPUT_PATH "bin" CACHE STRING "Path to place executables relative to ${CMAKE_INSTALL_PREFIX}")
+set(LIBRARY_OUTPUT_PATH "bin" CACHE STRING "Path to place libraries relative to ${CMAKE_INSTALL_PREFIX}")
target_link_libraries(xmr-stak ${LIBS} xmr-stak-c xmr-stak-backend)
@@ -535,23 +540,23 @@ target_link_libraries(xmr-stak ${LIBS} xmr-stak-c xmr-stak-backend)
# do not install the binary if the project and install are equal
if( NOT CMAKE_INSTALL_PREFIX STREQUAL PROJECT_BINARY_DIR )
install(TARGETS xmr-stak
- RUNTIME DESTINATION "${CMAKE_INSTALL_PREFIX}/bin")
+ RUNTIME DESTINATION "${CMAKE_INSTALL_PREFIX}/${EXECUTABLE_OUTPUT_PATH}")
if(CUDA_FOUND)
if(WIN32)
install(TARGETS xmrstak_cuda_backend
- RUNTIME DESTINATION "${CMAKE_INSTALL_PREFIX}/bin")
+ RUNTIME DESTINATION "${CMAKE_INSTALL_PREFIX}/${LIBRARY_OUTPUT_PATH}")
else()
install(TARGETS xmrstak_cuda_backend
- LIBRARY DESTINATION "${CMAKE_INSTALL_PREFIX}/bin")
+ LIBRARY DESTINATION "${CMAKE_INSTALL_PREFIX}/${LIBRARY_OUTPUT_PATH}")
endif()
endif()
if(OpenCL_FOUND)
if(WIN32)
install(TARGETS xmrstak_opencl_backend
- RUNTIME DESTINATION "${CMAKE_INSTALL_PREFIX}/bin")
+ RUNTIME DESTINATION "${CMAKE_INSTALL_PREFIX}/${LIBRARY_OUTPUT_PATH}")
else()
install(TARGETS xmrstak_opencl_backend
- LIBRARY DESTINATION "${CMAKE_INSTALL_PREFIX}/bin")
+ LIBRARY DESTINATION "${CMAKE_INSTALL_PREFIX}/${LIBRARY_OUTPUT_PATH}")
endif()
endif()
else()
diff --git a/README.md b/README.md
index 047aa510fdf783195df239f6c62b60662904b043..64fd4881b5e5b8d482f0bf981ccc22d0d103f947 100644
--- a/README.md
+++ b/README.md
@@ -41,9 +41,12 @@ XMR-Stak is a universal Stratum pool miner. This miner supports CPUs, AMD and NV
Besides [Monero](https://getmonero.org), following coins can be mined using this miner:
- [Aeon](http://www.aeon.cash)
+- [BBSCoin](https://www.bbscoin.xyz)
+- [Croat](https://croat.cat)
- [Edollar](https://edollar.cash)
- [Electroneum](https://electroneum.com)
- [Graft](https://www.graft.network)
+- [Haven](https://havenprotocol.com)
- [Intense](https://intensecoin.com)
- [Karbo](https://karbo.io)
- [Sumokoin](https://www.sumokoin.org)
diff --git a/doc/FAQ.md b/doc/FAQ.md
index 35dc2488dfa1ce5db6a7c08277b063405bb40caa..0661073205dea69d3da6ebd6c08c3312548bf0ea 100644
--- a/doc/FAQ.md
+++ b/doc/FAQ.md
@@ -9,7 +9,7 @@
* [Virus Protection Alert](#virus-protection-alert)
* [Change Currency to Mine](#change-currency-to-mine)
* [How can I mine Monero](#how-can-i-mine-monero)
-* [Why is Monero named monero7](why-is-monero-named-monero7)
+* [Why is Monero named monero7](#why-is-monero-named-monero7)
* [Which currency must be chosen if my fork coin is not listed](#which-currency-must-be-chosen-if-my-fork-coin-is-not-listed)
## "Obtaining SeLockMemoryPrivilege failed."
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp
index 006a7ed9e5680bcff28cb024c8efa56341ab148e..03100d0a3bbcff1003d433c84daf0eb2c97a8c6f 100644
--- a/xmrstak/backend/amd/amd_gpu/gpu.cpp
+++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp
@@ -26,6 +26,7 @@
#include <algorithm>
#include <regex>
#include <cassert>
+#include <algorithm>
#include <fstream>
#include <sstream>
@@ -307,12 +308,13 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_
return ERR_OCL_API;
}
- size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
- int threadMemMask = cn_select_mask(::jconf::inst()->GetMiningAlgo());
- int hashIterations = cn_select_iter(::jconf::inst()->GetMiningAlgo());
+ size_t scratchPadSize = std::max(
+ cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()),
+ cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot())
+ );
size_t g_thd = ctx->rawIntensity;
- ctx->ExtraBuffers[0] = clCreateBuffer(opencl_ctx, CL_MEM_READ_WRITE, hashMemSize * g_thd, NULL, &ret);
+ ctx->ExtraBuffers[0] = clCreateBuffer(opencl_ctx, CL_MEM_READ_WRITE, scratchPadSize * g_thd, NULL, &ret);
if(ret != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clCreateBuffer to create hash scratchpads buffer.", err_to_str(ret));
@@ -373,167 +375,203 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_
return ERR_OCL_API;
}
- auto miner_algo = ::jconf::inst()->GetMiningAlgo();
-
- char options[512];
- snprintf(options, sizeof(options),
- "-DITERATIONS=%d -DMASK=%d -DWORKSIZE=%llu -DSTRIDED_INDEX=%d -DMEM_CHUNK_EXPONENT=%d -DCOMP_MODE=%d -DMEMORY=%llu -DALGO=%d",
- hashIterations, threadMemMask, int_port(ctx->workSize), ctx->stridedIndex, int(1u<<ctx->memChunk), ctx->compMode ? 1 : 0,
- int_port(hashMemSize), int(miner_algo));
- /* create a hash for the compile time cache
- * used data:
- * - source code
- * - device name
- * - compile paramater
- */
- std::string src_str(source_code);
- src_str += options;
- src_str += devNameVec.data();
- std::string hash_hex_str;
- picosha2::hash256_hex_string(src_str, hash_hex_str);
-
- std::string cache_file = get_home() + "/.openclcache/" + hash_hex_str + ".openclbin";
- std::ifstream clBinFile(cache_file, std::ofstream::in | std::ofstream::binary);
- if(xmrstak::params::inst().AMDCache == false || !clBinFile.good())
- {
- if(xmrstak::params::inst().AMDCache)
- printer::inst()->print_msg(L1,"OpenCL device %u - Precompiled code %s not found. Compiling ...",ctx->deviceIdx, cache_file.c_str());
- ctx->Program = clCreateProgramWithSource(opencl_ctx, 1, (const char**)&source_code, NULL, &ret);
- if(ret != CL_SUCCESS)
- {
- printer::inst()->print_msg(L1,"Error %s when calling clCreateProgramWithSource on the OpenCL miner code", err_to_str(ret));
- return ERR_OCL_API;
- }
-
- ret = clBuildProgram(ctx->Program, 1, &ctx->DeviceID, options, NULL, NULL);
- if(ret != CL_SUCCESS)
+ xmrstak_algo miner_algo[2] = {
+ ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo(),
+ ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot()
+ };
+ int num_algos = miner_algo[0] == miner_algo[1] ? 1 : 2;
+
+ for(int ii = 0; ii < num_algos; ++ii)
+ {
+ // scratchpad size for the selected mining algorithm
+ size_t hashMemSize = cn_select_memory(miner_algo[ii]);
+ int threadMemMask = cn_select_mask(miner_algo[ii]);
+ int hashIterations = cn_select_iter(miner_algo[ii]);
+
+ char options[512];
+ snprintf(options, sizeof(options),
+ "-DITERATIONS=%d -DMASK=%d -DWORKSIZE=%llu -DSTRIDED_INDEX=%d -DMEM_CHUNK_EXPONENT=%d -DCOMP_MODE=%d -DMEMORY=%llu -DALGO=%d",
+ hashIterations, threadMemMask, int_port(ctx->workSize), ctx->stridedIndex, int(1u<<ctx->memChunk), ctx->compMode ? 1 : 0,
+ int_port(hashMemSize), int(miner_algo[ii]));
+ /* create a hash for the compile time cache
+ * used data:
+ * - source code
+ * - device name
+ * - compile paramater
+ */
+ std::string src_str(source_code);
+ src_str += options;
+ src_str += devNameVec.data();
+ std::string hash_hex_str;
+ picosha2::hash256_hex_string(src_str, hash_hex_str);
+
+ std::string cache_file = get_home() + "/.openclcache/" + hash_hex_str + ".openclbin";
+ std::ifstream clBinFile(cache_file, std::ofstream::in | std::ofstream::binary);
+ if(xmrstak::params::inst().AMDCache == false || !clBinFile.good())
{
- size_t len;
- printer::inst()->print_msg(L1,"Error %s when calling clBuildProgram.", err_to_str(ret));
-
- if((ret = clGetProgramBuildInfo(ctx->Program, ctx->DeviceID, CL_PROGRAM_BUILD_LOG, 0, NULL, &len)) != CL_SUCCESS)
+ if(xmrstak::params::inst().AMDCache)
+ printer::inst()->print_msg(L1,"OpenCL device %u - Precompiled code %s not found. Compiling ...",ctx->deviceIdx, cache_file.c_str());
+ ctx->Program[ii] = clCreateProgramWithSource(opencl_ctx, 1, (const char**)&source_code, NULL, &ret);
+ if(ret != CL_SUCCESS)
{
- printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for length of build log output.", err_to_str(ret));
+ printer::inst()->print_msg(L1,"Error %s when calling clCreateProgramWithSource on the OpenCL miner code", err_to_str(ret));
return ERR_OCL_API;
}
- char* BuildLog = (char*)malloc(len + 1);
- BuildLog[0] = '\0';
-
- if((ret = clGetProgramBuildInfo(ctx->Program, ctx->DeviceID, CL_PROGRAM_BUILD_LOG, len, BuildLog, NULL)) != CL_SUCCESS)
+ ret = clBuildProgram(ctx->Program[ii], 1, &ctx->DeviceID, options, NULL, NULL);
+ if(ret != CL_SUCCESS)
{
+ size_t len;
+ printer::inst()->print_msg(L1,"Error %s when calling clBuildProgram.", err_to_str(ret));
+
+ if((ret = clGetProgramBuildInfo(ctx->Program[ii], ctx->DeviceID, CL_PROGRAM_BUILD_LOG, 0, NULL, &len)) != CL_SUCCESS)
+ {
+ printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for length of build log output.", err_to_str(ret));
+ return ERR_OCL_API;
+ }
+
+ char* BuildLog = (char*)malloc(len + 1);
+ BuildLog[0] = '\0';
+
+ if((ret = clGetProgramBuildInfo(ctx->Program[ii], ctx->DeviceID, CL_PROGRAM_BUILD_LOG, len, BuildLog, NULL)) != CL_SUCCESS)
+ {
+ free(BuildLog);
+ printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for build log.", err_to_str(ret));
+ return ERR_OCL_API;
+ }
+
+ printer::inst()->print_str("Build log:\n");
+ std::cerr<<BuildLog<<std::endl;
+
free(BuildLog);
- printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for build log.", err_to_str(ret));
return ERR_OCL_API;
}
- printer::inst()->print_str("Build log:\n");
- std::cerr<<BuildLog<<std::endl;
+ cl_uint num_devices;
+ clGetProgramInfo(ctx->Program[ii], CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint), &num_devices,NULL);
- free(BuildLog);
- return ERR_OCL_API;
- }
- cl_uint num_devices;
- clGetProgramInfo(ctx->Program, CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint), &num_devices,NULL);
-
-
- std::vector<cl_device_id> devices_ids(num_devices);
- clGetProgramInfo(ctx->Program, CL_PROGRAM_DEVICES, sizeof(cl_device_id)* devices_ids.size(), devices_ids.data(),NULL);
- int dev_id = 0;
- /* Search for the gpu within the program context.
- * The id can be different to ctx->DeviceID.
- */
- for(auto & ocl_device : devices_ids)
- {
- if(ocl_device == ctx->DeviceID)
- break;
- dev_id++;
- }
+ std::vector<cl_device_id> devices_ids(num_devices);
+ clGetProgramInfo(ctx->Program[ii], CL_PROGRAM_DEVICES, sizeof(cl_device_id)* devices_ids.size(), devices_ids.data(),NULL);
+ int dev_id = 0;
+ /* Search for the gpu within the program context.
+ * The id can be different to ctx->DeviceID.
+ */
+ for(auto & ocl_device : devices_ids)
+ {
+ if(ocl_device == ctx->DeviceID)
+ break;
+ dev_id++;
+ }
- cl_build_status status;
- do
- {
- if((ret = clGetProgramBuildInfo(ctx->Program, ctx->DeviceID, CL_PROGRAM_BUILD_STATUS, sizeof(cl_build_status), &status, NULL)) != CL_SUCCESS)
+ cl_build_status status;
+ do
{
- printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for status of build.", err_to_str(ret));
- return ERR_OCL_API;
+ if((ret = clGetProgramBuildInfo(ctx->Program[ii], ctx->DeviceID, CL_PROGRAM_BUILD_STATUS, sizeof(cl_build_status), &status, NULL)) != CL_SUCCESS)
+ {
+ printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for status of build.", err_to_str(ret));
+ return ERR_OCL_API;
+ }
+ port_sleep(1);
}
- port_sleep(1);
- }
- while(status == CL_BUILD_IN_PROGRESS);
+ while(status == CL_BUILD_IN_PROGRESS);
- std::vector<size_t> binary_sizes(num_devices);
- clGetProgramInfo (ctx->Program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t) * binary_sizes.size(), binary_sizes.data(), NULL);
+ std::vector<size_t> binary_sizes(num_devices);
+ clGetProgramInfo (ctx->Program[ii], CL_PROGRAM_BINARY_SIZES, sizeof(size_t) * binary_sizes.size(), binary_sizes.data(), NULL);
- std::vector<char*> all_programs(num_devices);
- std::vector<std::vector<char>> program_storage;
+ std::vector<char*> all_programs(num_devices);
+ std::vector<std::vector<char>> program_storage;
+
+ if(xmrstak::params::inst().AMDCache)
+ {
+ int p_id = 0;
+ size_t mem_size = 0;
+ // create memory structure to query all OpenCL program binaries
+ for(auto & p : all_programs)
+ {
+ program_storage.emplace_back(std::vector<char>(binary_sizes[p_id]));
+ all_programs[p_id] = program_storage[p_id].data();
+ mem_size += binary_sizes[p_id];
+ p_id++;
+ }
+
+ if((ret = clGetProgramInfo(ctx->Program[ii], CL_PROGRAM_BINARIES, num_devices * sizeof(char*), all_programs.data(),NULL)) != CL_SUCCESS)
+ {
+ printer::inst()->print_msg(L1,"Error %s when calling clGetProgramInfo.", err_to_str(ret));
+ return ERR_OCL_API;
+ }
- if(xmrstak::params::inst().AMDCache)
+ std::ofstream file_stream;
+ file_stream.open(cache_file, std::ofstream::out | std::ofstream::binary);
+ file_stream.write(all_programs[dev_id], binary_sizes[dev_id]);
+ file_stream.close();
+ printer::inst()->print_msg(L1, "OpenCL device %u - Precompiled code stored in file %s",ctx->deviceIdx, cache_file.c_str());
+ }
+ }
+ else
{
- int p_id = 0;
- size_t mem_size = 0;
- // create memory structure to query all OpenCL program binaries
- for(auto & p : all_programs)
+ printer::inst()->print_msg(L1, "OpenCL device %u - Load precompiled code from file %s",ctx->deviceIdx, cache_file.c_str());
+ std::ostringstream ss;
+ ss << clBinFile.rdbuf();
+ std::string s = ss.str();
+
+ size_t bin_size = s.size();
+ auto data_ptr = s.data();
+
+ cl_int clStatus;
+ ctx->Program[ii] = clCreateProgramWithBinary(
+ opencl_ctx, 1, &ctx->DeviceID, &bin_size,
+ (const unsigned char **)&data_ptr, &clStatus, &ret
+ );
+ if(ret != CL_SUCCESS)
{
- program_storage.emplace_back(std::vector<char>(binary_sizes[p_id]));
- all_programs[p_id] = program_storage[p_id].data();
- mem_size += binary_sizes[p_id];
- p_id++;
+ printer::inst()->print_msg(L1,"Error %s when calling clCreateProgramWithBinary. Try to delete file %s", err_to_str(ret), cache_file.c_str());
+ return ERR_OCL_API;
}
-
- if((ret = clGetProgramInfo(ctx->Program, CL_PROGRAM_BINARIES, num_devices * sizeof(char*), all_programs.data(),NULL)) != CL_SUCCESS)
+ ret = clBuildProgram(ctx->Program[ii], 1, &ctx->DeviceID, NULL, NULL, NULL);
+ if(ret != CL_SUCCESS)
{
- printer::inst()->print_msg(L1,"Error %s when calling clGetProgramInfo.", err_to_str(ret));
+ printer::inst()->print_msg(L1,"Error %s when calling clBuildProgram. Try to delete file %s", err_to_str(ret), cache_file.c_str());
return ERR_OCL_API;
}
-
- std::ofstream file_stream;
- file_stream.open(cache_file, std::ofstream::out | std::ofstream::binary);
- file_stream.write(all_programs[dev_id], binary_sizes[dev_id]);
- file_stream.close();
- printer::inst()->print_msg(L1, "OpenCL device %u - Precompiled code stored in file %s",ctx->deviceIdx, cache_file.c_str());
}
- }
- else
- {
- printer::inst()->print_msg(L1, "OpenCL device %u - Load precompiled cod from file %s",ctx->deviceIdx, cache_file.c_str());
- std::ostringstream ss;
- ss << clBinFile.rdbuf();
- std::string s = ss.str();
-
- size_t bin_size = s.size();
- auto data_ptr = s.data();
-
- cl_int clStatus;
- ctx->Program = clCreateProgramWithBinary(
- opencl_ctx, 1, &ctx->DeviceID, &bin_size,
- (const unsigned char **)&data_ptr, &clStatus, &ret
- );
- if(ret != CL_SUCCESS)
+
+ std::vector<std::string> KernelNames = { "cn0", "cn1", "cn2", "Blake", "Groestl", "JH", "Skein" };
+ // append algorithm number to kernel name
+ for(int k = 0; k < 3; k++)
+ KernelNames[k] += std::to_string(miner_algo[ii]);
+
+ if(ii == 0)
{
- printer::inst()->print_msg(L1,"Error %s when calling clCreateProgramWithBinary. Try to delete file %s", err_to_str(ret), cache_file.c_str());
- return ERR_OCL_API;
+ for(int i = 0; i < 7; ++i)
+ {
+ ctx->Kernels[ii][i] = clCreateKernel(ctx->Program[ii], KernelNames[i].c_str(), &ret);
+ if(ret != CL_SUCCESS)
+ {
+ printer::inst()->print_msg(L1,"Error %s when calling clCreateKernel for kernel_0 %s.", err_to_str(ret), KernelNames[i].c_str());
+ return ERR_OCL_API;
+ }
+ }
}
- ret = clBuildProgram(ctx->Program, 1, &ctx->DeviceID, NULL, NULL, NULL);
- if(ret != CL_SUCCESS)
+ else
{
- printer::inst()->print_msg(L1,"Error %s when calling clBuildProgram. Try to delete file %s", err_to_str(ret), cache_file.c_str());
- return ERR_OCL_API;
- }
- }
+ for(int i = 0; i < 3; ++i)
+ {
+ ctx->Kernels[ii][i] = clCreateKernel(ctx->Program[ii], KernelNames[i].c_str(), &ret);
+ if(ret != CL_SUCCESS)
+ {
+ printer::inst()->print_msg(L1,"Error %s when calling clCreateKernel for kernel_1 %s.", err_to_str(ret), KernelNames[i].c_str());
+ return ERR_OCL_API;
+ }
+ }
+ // move kernel from the main algorithm into the root algorithm kernel space
+ for(int i = 3; i < 7; ++i)
+ {
+ ctx->Kernels[ii][i] = ctx->Kernels[0][i];
+ }
- const char *KernelNames[] = { "cn0", "cn1", "cn2", "Blake", "Groestl", "JH", "Skein", "cn1_monero" };
- for(int i = 0; i < 8; ++i)
- {
- ctx->Kernels[i] = clCreateKernel(ctx->Program, KernelNames[i], &ret);
- if(ret != CL_SUCCESS)
- {
- printer::inst()->print_msg(L1,"Error %s when calling clCreateKernel for kernel %s.", err_to_str(ret), KernelNames[i]);
- return ERR_OCL_API;
}
}
-
ctx->Nonce = 0;
return 0;
}
@@ -637,11 +675,18 @@ std::vector<GpuContext> getAMDDevices(int index)
}
std::string devVendor(devVendorVec.data());
- if( devVendor.find("Advanced Micro Devices") != std::string::npos || devVendor.find("AMD") != std::string::npos)
+
+ bool isAMDDevice = devVendor.find("Advanced Micro Devices") != std::string::npos || devVendor.find("AMD") != std::string::npos;
+ bool isNVIDIADevice = devVendor.find("NVIDIA Corporation") != std::string::npos || devVendor.find("NVIDIA") != std::string::npos;
+
+ std::string selectedOpenCLVendor = xmrstak::params::inst().openCLVendor;
+ if((isAMDDevice && selectedOpenCLVendor == "AMD") || (isNVIDIADevice && selectedOpenCLVendor == "NVIDIA"))
{
GpuContext ctx;
std::vector<char> devNameVec(1024);
size_t maxMem;
+ if( devVendor.find("NVIDIA Corporation") != std::string::npos)
+ ctx.isNVIDIA = true;
if((clStatus = clGetDeviceInfo(device_list[k], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(int), &(ctx.computeUnits), NULL)) != CL_SUCCESS)
{
@@ -661,6 +706,10 @@ std::vector<GpuContext> getAMDDevices(int index)
continue;
}
+ // the allocation for NVIDIA OpenCL is not limited to 1/4 of the GPU memory per allocation
+ if(ctx.isNVIDIA)
+ maxMem = ctx.freeMem;
+
if((clStatus = clGetDeviceInfo(device_list[k], CL_DEVICE_NAME, devNameVec.size(), devNameVec.data(), NULL)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"WARNING: %s when calling clGetDeviceInfo to get CL_DEVICE_NAME for device %u.", err_to_str(clStatus), k);
@@ -709,13 +758,15 @@ int getAMDPlatformIdx()
clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, infoSize, platformNameVec.data(), NULL);
std::string platformName(platformNameVec.data());
- if( platformName.find("Advanced Micro Devices") != std::string::npos ||
+
+ bool isAMDOpenCL = platformName.find("Advanced Micro Devices") != std::string::npos ||
platformName.find("Apple") != std::string::npos ||
- platformName.find("Mesa") != std::string::npos
- )
+ platformName.find("Mesa") != std::string::npos;
+ bool isNVIDIADevice = platformName.find("NVIDIA Corporation") != std::string::npos || platformName.find("NVIDIA") != std::string::npos;
+ std::string selectedOpenCLVendor = xmrstak::params::inst().openCLVendor;
+ if((isAMDOpenCL && selectedOpenCLVendor == "AMD") || (isNVIDIADevice && selectedOpenCLVendor == "NVIDIA"))
{
-
- printer::inst()->print_msg(L0,"Found AMD platform index id = %i, name = %s",i , platformName.c_str());
+ printer::inst()->print_msg(L0,"Found %s platform index id = %i, name = %s", selectedOpenCLVendor.c_str(), i , platformName.c_str());
if(platformName.find("Mesa") != std::string::npos)
mesaPlatform = i;
else
@@ -781,7 +832,7 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx)
std::vector<char> platformNameVec(infoSize);
clGetPlatformInfo(PlatformIDList[platform_idx], CL_PLATFORM_VENDOR, infoSize, platformNameVec.data(), NULL);
std::string platformName(platformNameVec.data());
- if( platformName.find("Advanced Micro Devices") == std::string::npos)
+ if(xmrstak::params::inst().openCLVendor == "AMD" && platformName.find("Advanced Micro Devices") == std::string::npos)
{
printer::inst()->print_msg(L1,"WARNING: using non AMD device: %s", platformName.c_str());
}
@@ -869,7 +920,8 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx)
{
size_t reduced_intensity = (ctx[i].rawIntensity / ctx[i].workSize) * ctx[i].workSize;
ctx[i].rawIntensity = reduced_intensity;
- printer::inst()->print_msg(L0, "WARNING AMD: gpu %d intensity is not a multiple of 'worksize', auto reduce intensity to %d", ctx[i].deviceIdx, int(reduced_intensity));
+ const std::string backendName = xmrstak::params::inst().openCLVendor;
+ printer::inst()->print_msg(L0, "WARNING %s: gpu %d intensity is not a multiple of 'worksize', auto reduce intensity to %d", backendName.c_str(), ctx[i].deviceIdx, int(reduced_intensity));
}
if((ret = InitOpenCLGpu(opencl_ctx, &ctx[i], source_code.c_str())) != ERR_SUCCESS)
@@ -881,8 +933,11 @@ 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, xmrstak_algo miner_algo, uint32_t version)
+size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target, xmrstak_algo miner_algo)
{
+ // switch to the kernel storage
+ int kernel_storage = miner_algo == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() ? 0 : 1;
+
cl_int ret;
if(input_len > 84)
@@ -899,71 +954,51 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
return ERR_OCL_API;
}
- if((ret = clSetKernelArg(ctx->Kernels[0], 0, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][0], 0, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 0, argument 0.", err_to_str(ret));
return ERR_OCL_API;
}
// Scratchpads
- if((ret = clSetKernelArg(ctx->Kernels[0], 1, sizeof(cl_mem), ctx->ExtraBuffers + 0)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][0], 1, sizeof(cl_mem), ctx->ExtraBuffers + 0)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 0, argument 1.", err_to_str(ret));
return ERR_OCL_API;
}
// States
- if((ret = clSetKernelArg(ctx->Kernels[0], 2, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][0], 2, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 0, argument 2.", err_to_str(ret));
return ERR_OCL_API;
}
// Threads
- if((ret = clSetKernelArg(ctx->Kernels[0], 3, sizeof(cl_ulong), &numThreads)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][0], 3, sizeof(cl_ulong), &numThreads)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 0, argument 3.", err_to_str(ret));
return(ERR_OCL_API);
}
- /* ATTENTION: if we miner cryptonight_heavy the kernel needs an additional parameter version.
- * Do NOT use the variable `miner_algo` because this variable is changed dynamicly
- */
- if(::jconf::inst()->GetMiningAlgo() == cryptonight_heavy)
- {
- // version
- if ((ret = clSetKernelArg(ctx->Kernels[0], 4, sizeof(cl_uint), &version)) != CL_SUCCESS)
- {
- printer::inst()->print_msg(L1, "Error %s when calling clSetKernelArg for kernel 0, argument 4.", err_to_str(ret));
- return ERR_OCL_API;
- }
- }
-
// CN1 Kernel
- /// @todo only activate if currency is monero
- int cn_kernel_offset = 0;
- if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon)
- {
- cn_kernel_offset = 6;
- }
-
// Scratchpads
- if((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 0, sizeof(cl_mem), ctx->ExtraBuffers + 0)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][1], 0, sizeof(cl_mem), ctx->ExtraBuffers + 0)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 1, argument 0.", err_to_str(ret));
return ERR_OCL_API;
}
// States
- if((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 1, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][1], 1, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 1, argument 1.", err_to_str(ret));
return ERR_OCL_API;
}
// Threads
- if((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 2, sizeof(cl_ulong), &numThreads)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][1], 2, sizeof(cl_ulong), &numThreads)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 1, argument 2.", err_to_str(ret));
return(ERR_OCL_API);
@@ -972,113 +1007,88 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon )
{
// Input
- if ((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS)
+ if ((ret = clSetKernelArg(ctx->Kernels[kernel_storage][1], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1, "Error %s when calling clSetKernelArg for kernel 1, arugment 4(input buffer).", err_to_str(ret));
return ERR_OCL_API;
}
}
- /* ATTENTION: if we miner cryptonight_heavy the kernel needs an additional parameter version.
- * Do NOT use the variable `miner_algo` because this variable is changed dynamicly
- */
- else if(::jconf::inst()->GetMiningAlgo() == cryptonight_heavy)
- {
- // version
- if ((ret = clSetKernelArg(ctx->Kernels[1], 3, sizeof(cl_uint), &version)) != CL_SUCCESS)
- {
- printer::inst()->print_msg(L1, "Error %s when calling clSetKernelArg for kernel 1, argument 3 (version).", err_to_str(ret));
- return ERR_OCL_API;
- }
- }
// CN3 Kernel
// Scratchpads
- if((ret = clSetKernelArg(ctx->Kernels[2], 0, sizeof(cl_mem), ctx->ExtraBuffers + 0)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][2], 0, sizeof(cl_mem), ctx->ExtraBuffers + 0)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 2, argument 0.", err_to_str(ret));
return ERR_OCL_API;
}
// States
- if((ret = clSetKernelArg(ctx->Kernels[2], 1, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][2], 1, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 2, argument 1.", err_to_str(ret));
return ERR_OCL_API;
}
// Branch 0
- if((ret = clSetKernelArg(ctx->Kernels[2], 2, sizeof(cl_mem), ctx->ExtraBuffers + 2)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][2], 2, sizeof(cl_mem), ctx->ExtraBuffers + 2)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 2, argument 2.", err_to_str(ret));
return ERR_OCL_API;
}
// Branch 1
- if((ret = clSetKernelArg(ctx->Kernels[2], 3, sizeof(cl_mem), ctx->ExtraBuffers + 3)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][2], 3, sizeof(cl_mem), ctx->ExtraBuffers + 3)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 2, argument 3.", err_to_str(ret));
return ERR_OCL_API;
}
// Branch 2
- if((ret = clSetKernelArg(ctx->Kernels[2], 4, sizeof(cl_mem), ctx->ExtraBuffers + 4)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][2], 4, sizeof(cl_mem), ctx->ExtraBuffers + 4)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 2, argument 4.", err_to_str(ret));
return ERR_OCL_API;
}
// Branch 3
- if((ret = clSetKernelArg(ctx->Kernels[2], 5, sizeof(cl_mem), ctx->ExtraBuffers + 5)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][2], 5, sizeof(cl_mem), ctx->ExtraBuffers + 5)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 2, argument 5.", err_to_str(ret));
return ERR_OCL_API;
}
// Threads
- if((ret = clSetKernelArg(ctx->Kernels[2], 6, sizeof(cl_ulong), &numThreads)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][2], 6, sizeof(cl_ulong), &numThreads)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 2, argument 6.", err_to_str(ret));
return(ERR_OCL_API);
}
- /* ATTENTION: if we miner cryptonight_heavy the kernel needs an additional parameter version.
- * Do NOT use the variable `miner_algo` because this variable is changed dynamicly
- */
- if(::jconf::inst()->GetMiningAlgo() == cryptonight_heavy)
- {
- // version
- if ((ret = clSetKernelArg(ctx->Kernels[2], 7, sizeof(cl_uint), &version)) != CL_SUCCESS)
- {
- printer::inst()->print_msg(L1, "Error %s when calling clSetKernelArg for kernel 2, argument 7.", err_to_str(ret));
- return ERR_OCL_API;
- }
- }
-
for(int i = 0; i < 4; ++i)
{
// States
- if((ret = clSetKernelArg(ctx->Kernels[i + 3], 0, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][i + 3], 0, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel %d, argument %d.", err_to_str(ret), i + 3, 0);
return ERR_OCL_API;
}
// Nonce buffer
- if((ret = clSetKernelArg(ctx->Kernels[i + 3], 1, sizeof(cl_mem), ctx->ExtraBuffers + (i + 2))) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][i + 3], 1, sizeof(cl_mem), ctx->ExtraBuffers + (i + 2))) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel %d, argument %d.", err_to_str(ret), i + 3, 1);
return ERR_OCL_API;
}
// Output
- if((ret = clSetKernelArg(ctx->Kernels[i + 3], 2, sizeof(cl_mem), &ctx->OutputBuffer)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][i + 3], 2, sizeof(cl_mem), &ctx->OutputBuffer)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel %d, argument %d.", err_to_str(ret), i + 3, 2);
return ERR_OCL_API;
}
// Target
- if((ret = clSetKernelArg(ctx->Kernels[i + 3], 3, sizeof(cl_ulong), &target)) != CL_SUCCESS)
+ if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][i + 3], 3, sizeof(cl_ulong), &target)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel %d, argument %d.", err_to_str(ret), i + 3, 3);
return ERR_OCL_API;
@@ -1088,8 +1098,11 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
return ERR_SUCCESS;
}
-size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo, uint32_t version)
+size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo)
{
+ // switch to the kernel storage
+ int kernel_storage = miner_algo == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() ? 0 : 1;
+
cl_int ret;
cl_uint zero = 0;
size_t BranchNonces[4];
@@ -1125,35 +1138,21 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo,
clFinish(ctx->CommandQueues);
size_t Nonce[2] = {ctx->Nonce, 1}, gthreads[2] = { g_thd, 8 }, lthreads[2] = { w_size, 8 };
- if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[0], 2, Nonce, gthreads, lthreads, 0, NULL, NULL)) != CL_SUCCESS)
+ if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[kernel_storage][0], 2, Nonce, gthreads, lthreads, 0, NULL, NULL)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clEnqueueNDRangeKernel for kernel %d.", err_to_str(ret), 0);
return ERR_OCL_API;
}
- /*for(int i = 1; i < 3; ++i)
- {
- if((ret = clEnqueueNDRangeKernel(*ctx->CommandQueues, ctx->Kernels[i], 1, &ctx->Nonce, &g_thd, &w_size, 0, NULL, NULL)) != CL_SUCCESS)
- {
- Log(LOG_CRITICAL, "Error %s when calling clEnqueueNDRangeKernel for kernel %d.", err_to_str(ret), i);
- return(ERR_OCL_API);
- }
- }*/
-
size_t tmpNonce = ctx->Nonce;
- /// @todo only activate if currency is monero
- int cn_kernel_offset = 0;
- if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon)
- {
- cn_kernel_offset = 6;
- }
- if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[1 + cn_kernel_offset], 1, &tmpNonce, &g_thd, &w_size, 0, NULL, NULL)) != CL_SUCCESS)
+
+ if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[kernel_storage][1], 1, &tmpNonce, &g_thd, &w_size, 0, NULL, NULL)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clEnqueueNDRangeKernel for kernel %d.", err_to_str(ret), 1);
return ERR_OCL_API;
}
- if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[2], 2, Nonce, gthreads, lthreads, 0, NULL, NULL)) != CL_SUCCESS)
+ if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[kernel_storage][2], 2, Nonce, gthreads, lthreads, 0, NULL, NULL)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clEnqueueNDRangeKernel for kernel %d.", err_to_str(ret), 2);
return ERR_OCL_API;
@@ -1190,7 +1189,7 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo,
if(BranchNonces[i])
{
// Threads
- if((clSetKernelArg(ctx->Kernels[i + 3], 4, sizeof(cl_ulong), BranchNonces + i)) != CL_SUCCESS)
+ if((clSetKernelArg(ctx->Kernels[kernel_storage][i + 3], 4, sizeof(cl_ulong), BranchNonces + i)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel %d, argument %d.", err_to_str(ret), i + 3, 4);
return(ERR_OCL_API);
@@ -1201,7 +1200,7 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo,
// number of global threads must be a multiple of the work group size (w_size)
assert(BranchNonces[i]%w_size == 0);
size_t tmpNonce = ctx->Nonce;
- if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[i + 3], 1, &tmpNonce, BranchNonces + i, &w_size, 0, NULL, NULL)) != CL_SUCCESS)
+ if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[kernel_storage][i + 3], 1, &tmpNonce, BranchNonces + i, &w_size, 0, NULL, NULL)) != CL_SUCCESS)
{
printer::inst()->print_msg(L1,"Error %s when calling clEnqueueNDRangeKernel for kernel %d.", err_to_str(ret), i + 3);
return ERR_OCL_API;
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.hpp b/xmrstak/backend/amd/amd_gpu/gpu.hpp
index a387b1535dd0b2c701e8d6043c7d7eac45b45eaf..5ab80b82a93d0f0c9fffee5b0846d2c8ecab0653 100644
--- a/xmrstak/backend/amd/amd_gpu/gpu.hpp
+++ b/xmrstak/backend/amd/amd_gpu/gpu.hpp
@@ -27,6 +27,7 @@ struct GpuContext
size_t workSize;
int stridedIndex;
int memChunk;
+ bool isNVIDIA = false;
int compMode;
/*Output vars*/
@@ -35,8 +36,8 @@ struct GpuContext
cl_mem InputBuffer;
cl_mem OutputBuffer;
cl_mem ExtraBuffers[6];
- cl_program Program;
- cl_kernel Kernels[8];
+ cl_program Program[2];
+ cl_kernel Kernels[2][8];
size_t freeMem;
int computeUnits;
std::string name;
@@ -50,7 +51,7 @@ 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, xmrstak_algo miner_algo, uint32_t version);
-size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo, uint32_t version);
+size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target, xmrstak_algo miner_algo);
+size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo);
diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
index 5d4e66ca9170a535f7e2fe58f71213b3d4737630..d2ae1a7ee2ce43b69be110833fb14332c3eb6d41 100644
--- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
+++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
@@ -433,15 +433,13 @@ inline ulong getIdx()
#endif
}
-#define mix_and_propagate(xin) (xin)[(get_local_id(1)) % 8][get_local_id(0)] ^ (xin)[(get_local_id(1) + 1) % 8][get_local_id(0)]
+#define mix_and_propagate(xin) (xin)[(get_local_id(1)) % 8][get_local_id(0)] ^ (xin)[(get_local_id(1) + 1) % 8][get_local_id(0)]
+
+#define JOIN_DO(x,y) x##y
+#define JOIN(x,y) JOIN_DO(x,y)
__attribute__((reqd_work_group_size(WORKSIZE, 8, 1)))
-__kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states, ulong Threads
-// cryptonight_heavy
-#if (ALGO == 4)
- , uint version
-#endif
-)
+__kernel void JOIN(cn0,ALGO)(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states, ulong Threads)
{
ulong State[25];
uint ExpandedKey1[40];
@@ -517,23 +515,20 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul
// cryptonight_heavy
#if (ALGO == 4)
- if(version >= 3)
- {
- __local uint4 xin[8][WORKSIZE];
+ __local uint4 xin[8][WORKSIZE];
- /* Also left over threads performe this loop.
- * The left over thread results will be ignored
- */
- for(size_t i=0; i < 16; i++)
- {
- #pragma unroll
- for(int j = 0; j < 10; ++j)
- text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey1)[j]);
- barrier(CLK_LOCAL_MEM_FENCE);
- xin[get_local_id(1)][get_local_id(0)] = text;
- barrier(CLK_LOCAL_MEM_FENCE);
- text = mix_and_propagate(xin);
- }
+ /* Also left over threads performe this loop.
+ * The left over thread results will be ignored
+ */
+ for(size_t i=0; i < 16; i++)
+ {
+ #pragma unroll
+ for(int j = 0; j < 10; ++j)
+ text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey1)[j]);
+ barrier(CLK_LOCAL_MEM_FENCE);
+ xin[get_local_id(1)][get_local_id(0)] = text;
+ barrier(CLK_LOCAL_MEM_FENCE);
+ text = mix_and_propagate(xin);
}
#endif
@@ -542,13 +537,9 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul
if(gIdx < Threads)
#endif
{
- int iterations = MEMORY >> 7;
-#if (ALGO == 4)
- if(version < 3)
- iterations >>= 1;
-#endif
+
#pragma unroll 2
- for(int i = 0; i < iterations; ++i)
+ for(int i = 0; i < (MEMORY >> 7); ++i)
{
#pragma unroll
for(int j = 0; j < 10; ++j)
@@ -560,22 +551,13 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul
mem_fence(CLK_GLOBAL_MEM_FENCE);
}
-#define VARIANT1_1(p) \
- uint table = 0x75310U; \
- uint index = (((p).s2 >> 26) & 12) | (((p).s2 >> 23) & 2); \
- (p).s2 ^= ((table >> index) & 0x30U) << 24
-
-#define VARIANT1_2(p) ((uint2 *)&(p))[0] ^= tweak1_2
-
-#define VARIANT1_INIT() \
- tweak1_2 = as_uint2(input[4]); \
- tweak1_2.s0 >>= 24; \
- tweak1_2.s0 |= tweak1_2.s1 << 8; \
- tweak1_2.s1 = get_global_id(0); \
- tweak1_2 ^= as_uint2(states[24])
-
__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
-__kernel void cn1_monero(__global uint4 *Scratchpad, __global ulong *states, ulong Threads, __global ulong *input)
+__kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states, ulong Threads
+// cryptonight_monero || cryptonight_aeon
+#if(ALGO == 3 || ALGO == 5)
+, __global ulong *input
+#endif
+)
{
ulong a[2], b[2];
__local uint AES0[256], AES1[256], AES2[256], AES3[256];
@@ -592,8 +574,9 @@ __kernel void cn1_monero(__global uint4 *Scratchpad, __global ulong *states, ulo
}
barrier(CLK_LOCAL_MEM_FENCE);
-
+#if(ALGO == 3 || ALGO == 5)
uint2 tweak1_2;
+#endif
uint4 b_x;
#if(COMP_MODE==1)
// do not use early return here
@@ -615,7 +598,13 @@ __kernel void cn1_monero(__global uint4 *Scratchpad, __global ulong *states, ulo
b[1] = states[3] ^ states[7];
b_x = ((uint4 *)b)[0];
- VARIANT1_INIT();
+#if(ALGO == 3 || ALGO == 5)
+ tweak1_2 = as_uint2(input[4]);
+ tweak1_2.s0 >>= 24;
+ tweak1_2.s0 |= tweak1_2.s1 << 8;
+ tweak1_2.s1 = get_global_id(0);
+ tweak1_2 ^= as_uint2(states[24]);
+#endif
}
mem_fence(CLK_LOCAL_MEM_FENCE);
@@ -625,17 +614,23 @@ __kernel void cn1_monero(__global uint4 *Scratchpad, __global ulong *states, ulo
if(gIdx < Threads)
#endif
{
+ ulong idx0 = a[0];
+
#pragma unroll 8
for(int i = 0; i < ITERATIONS; ++i)
{
ulong c[2];
- ((uint4 *)c)[0] = Scratchpad[IDX((a[0] & MASK) >> 4)];
+ ((uint4 *)c)[0] = Scratchpad[IDX((idx0 & MASK) >> 4)];
((uint4 *)c)[0] = AES_Round(AES0, AES1, AES2, AES3, ((uint4 *)c)[0], ((uint4 *)a)[0]);
b_x ^= ((uint4 *)c)[0];
- VARIANT1_1(b_x);
- Scratchpad[IDX((a[0] & MASK) >> 4)] = b_x;
+#if(ALGO == 3 || ALGO == 5)
+ uint table = 0x75310U;
+ uint index = ((b_x.s2 >> 26) & 12) | ((b_x.s2 >> 23) & 2);
+ b_x.s2 ^= ((table >> index) & 0x30U) << 24;
+#endif
+ Scratchpad[IDX((idx0 & MASK) >> 4)] = b_x;
uint4 tmp;
tmp = Scratchpad[IDX((c[0] & MASK) >> 4)];
@@ -643,101 +638,14 @@ __kernel void cn1_monero(__global uint4 *Scratchpad, __global ulong *states, ulo
a[1] += c[0] * as_ulong2(tmp).s0;
a[0] += mul_hi(c[0], as_ulong2(tmp).s0);
- VARIANT1_2(a[1]);
- Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0];
- VARIANT1_2(a[1]);
-
- ((uint4 *)a)[0] ^= tmp;
-
- b_x = ((uint4 *)c)[0];
- }
- }
- mem_fence(CLK_GLOBAL_MEM_FENCE);
-}
-
-__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
-__kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Threads
-// cryptonight_heavy
-#if (ALGO == 4)
- , uint version
-#endif
-)
-{
- ulong a[2], b[2];
- __local uint AES0[256], AES1[256], AES2[256], AES3[256];
-
- const ulong gIdx = getIdx();
-
- 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);
-
- uint4 b_x;
-#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);
-#endif
-
- a[0] = states[0] ^ states[4];
- b[0] = states[2] ^ states[6];
- a[1] = states[1] ^ states[5];
- b[1] = states[3] ^ states[7];
-
- b_x = ((uint4 *)b)[0];
- }
-
- mem_fence(CLK_LOCAL_MEM_FENCE);
-
-#if(COMP_MODE==1)
- // do not use early return here
- if(gIdx < Threads)
-#endif
- {
- ulong idx0 = a[0];
- ulong mask = MASK;
- int iterations = ITERATIONS;
-#if (ALGO == 4)
- if(version < 3)
- {
- iterations <<= 1;
- mask -= 0x200000;
- }
+#if(ALGO == 3 || ALGO == 5)
+ ((uint2 *)&(a[1]))[0] ^= tweak1_2;
+ Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0];
+ ((uint2 *)&(a[1]))[0] ^= tweak1_2;
+#else
+ Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0];
#endif
- #pragma unroll 8
- for(int i = 0; i < iterations; ++i)
- {
- ulong c[2];
-
- ((uint4 *)c)[0] = Scratchpad[IDX((idx0 & mask) >> 4)];
- ((uint4 *)c)[0] = AES_Round(AES0, AES1, AES2, AES3, ((uint4 *)c)[0], ((uint4 *)a)[0]);
- //b_x ^= ((uint4 *)c)[0];
-
- Scratchpad[IDX((idx0 & mask) >> 4)] = b_x ^ ((uint4 *)c)[0];
-
- uint4 tmp;
- tmp = Scratchpad[IDX((c[0] & mask) >> 4)];
-
- a[1] += c[0] * as_ulong2(tmp).s0;
- a[0] += mul_hi(c[0], as_ulong2(tmp).s0);
-
- Scratchpad[IDX((c[0] & mask) >> 4)] = ((uint4 *)a)[0];
((uint4 *)a)[0] ^= tmp;
idx0 = a[0];
@@ -745,14 +653,11 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre
b_x = ((uint4 *)c)[0];
// cryptonight_heavy
#if (ALGO == 4)
- if(version >= 3)
- {
- long n = *((__global long*)(Scratchpad + (IDX((idx0 & mask) >> 4))));
- int d = ((__global int*)(Scratchpad + (IDX((idx0 & mask) >> 4))))[2];
- long q = n / (d | 0x5);
- *((__global long*)(Scratchpad + (IDX((idx0 & mask) >> 4)))) = n ^ q;
- idx0 = d ^ q;
- }
+ long n = *((__global long*)(Scratchpad + (IDX((idx0 & MASK) >> 4))));
+ int d = ((__global int*)(Scratchpad + (IDX((idx0 & MASK) >> 4))))[2];
+ long q = n / (d | 0x5);
+ *((__global long*)(Scratchpad + (IDX((idx0 & MASK) >> 4)))) = n ^ q;
+ idx0 = d ^ q;
#endif
}
}
@@ -760,12 +665,7 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre
}
__attribute__((reqd_work_group_size(WORKSIZE, 8, 1)))
-__kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global uint *Branch0, __global uint *Branch1, __global uint *Branch2, __global uint *Branch3, ulong Threads
-// cryptonight_heavy
-#if (ALGO == 4)
- , uint version
-#endif
- )
+__kernel void JOIN(cn2,ALGO) (__global uint4 *Scratchpad, __global ulong *states, __global uint *Branch0, __global uint *Branch1, __global uint *Branch2, __global uint *Branch3, ulong Threads)
{
__local uint AES0[256], AES1[256], AES2[256], AES3[256];
uint ExpandedKey2[40];
@@ -827,58 +727,42 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u
if(gIdx < Threads)
#endif
{
- int iterations = MEMORY >> 7;
#if (ALGO == 4)
- if(version < 3)
- {
- iterations >>= 1;
- #pragma unroll 2
- for(int i = 0; i < iterations; ++i)
- {
- text ^= Scratchpad[IDX((i << 3) + get_local_id(1))];
-
- #pragma unroll 10
- for(int j = 0; j < 10; ++j)
- text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]);
- }
- }
- else
+ #pragma unroll 2
+ for(int i = 0; i < (MEMORY >> 7); ++i)
{
- #pragma unroll 2
- for(int i = 0; i < iterations; ++i)
- {
- text ^= Scratchpad[IDX((i << 3) + get_local_id(1))];
+ text ^= Scratchpad[IDX((i << 3) + get_local_id(1))];
- #pragma unroll 10
- for(int j = 0; j < 10; ++j)
- text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]);
+ #pragma unroll 10
+ for(int j = 0; j < 10; ++j)
+ text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]);
- barrier(CLK_LOCAL_MEM_FENCE);
- xin[get_local_id(1)][get_local_id(0)] = text;
- barrier(CLK_LOCAL_MEM_FENCE);
- text = mix_and_propagate(xin);
- }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ xin[get_local_id(1)][get_local_id(0)] = text;
+ barrier(CLK_LOCAL_MEM_FENCE);
+ text = mix_and_propagate(xin);
+ }
- #pragma unroll 2
- for(int i = 0; i < iterations; ++i)
- {
- text ^= Scratchpad[IDX((i << 3) + get_local_id(1))];
+ #pragma unroll 2
+ for(int i = 0; i < (MEMORY >> 7); ++i)
+ {
+ text ^= Scratchpad[IDX((i << 3) + get_local_id(1))];
- #pragma unroll 10
- for(int j = 0; j < 10; ++j)
- text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]);
+ #pragma unroll 10
+ for(int j = 0; j < 10; ++j)
+ text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]);
- barrier(CLK_LOCAL_MEM_FENCE);
- xin[get_local_id(1)][get_local_id(0)] = text;
- barrier(CLK_LOCAL_MEM_FENCE);
- text = mix_and_propagate(xin);
- }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ xin[get_local_id(1)][get_local_id(0)] = text;
+ barrier(CLK_LOCAL_MEM_FENCE);
+ text = mix_and_propagate(xin);
}
+
#else
#pragma unroll 2
- for(int i = 0; i < iterations; ++i)
+ for(int i = 0; i < (MEMORY >> 7); ++i)
{
text ^= Scratchpad[IDX((i << 3) + get_local_id(1))];
@@ -891,21 +775,18 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u
// cryptonight_heavy
#if (ALGO == 4)
- if(version >= 3)
+ /* Also left over threads perform this loop.
+ * The left over thread results will be ignored
+ */
+ for(size_t i=0; i < 16; i++)
{
- /* Also left over threads performe this loop.
- * The left over thread results will be ignored
- */
- for(size_t i=0; i < 16; i++)
- {
- #pragma unroll
- for(int j = 0; j < 10; ++j)
- text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]);
- barrier(CLK_LOCAL_MEM_FENCE);
- xin[get_local_id(1)][get_local_id(0)] = text;
- barrier(CLK_LOCAL_MEM_FENCE);
- text = mix_and_propagate(xin);
- }
+ #pragma unroll
+ for(int j = 0; j < 10; ++j)
+ text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]);
+ barrier(CLK_LOCAL_MEM_FENCE);
+ xin[get_local_id(1)][get_local_id(0)] = text;
+ barrier(CLK_LOCAL_MEM_FENCE);
+ text = mix_and_propagate(xin);
}
#endif
diff --git a/xmrstak/backend/amd/autoAdjust.hpp b/xmrstak/backend/amd/autoAdjust.hpp
index ea057a0f6f1e8e50a1d5681dd81625bc4def7a26..685890bc27eafcb172fa941a1fc34575b6edbff5 100644
--- a/xmrstak/backend/amd/autoAdjust.hpp
+++ b/xmrstak/backend/amd/autoAdjust.hpp
@@ -83,11 +83,15 @@ private:
constexpr size_t byteToMiB = 1024u * 1024u;
- size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
+ size_t hashMemSize = std::max(
+ cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()),
+ cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot())
+ );
std::string conf;
for(auto& ctx : devVec)
{
+ size_t minFreeMem = 128u * byteToMiB;
/* 1000 is a magic selected limit, the reason is that more than 2GiB memory
* sowing down the memory performance because of TLB cache misses
*/
@@ -109,12 +113,26 @@ private:
*/
maxThreads = 2024u;
}
+
+ // NVIDIA optimizations
+ if(
+ ctx.isNVIDIA && (
+ ctx.name.find("P100") != std::string::npos ||
+ ctx.name.find("V100") != std::string::npos
+ )
+ )
+ {
+ // do not limit the number of threads
+ maxThreads = 40000u;
+ minFreeMem = 512u * byteToMiB;
+ }
+
// increase all intensity limits by two for aeon
- if(::jconf::inst()->GetMiningAlgo() == cryptonight_lite)
+ if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_lite)
maxThreads *= 2u;
// keep 128MiB memory free (value is randomly chosen)
- size_t availableMem = ctx.freeMem - (128u * byteToMiB);
+ size_t availableMem = ctx.freeMem - minFreeMem;
// 224byte extra memory is used per thread for meta data
size_t perThread = hashMemSize + 224u;
size_t maxIntensity = availableMem / perThread;
@@ -135,7 +153,7 @@ private:
// set 8 threads per block (this is a good value for the most gpus)
conf += std::string(" { \"index\" : ") + std::to_string(ctx.deviceIdx) + ",\n" +
" \"intensity\" : " + std::to_string(intensity) + ", \"worksize\" : " + std::to_string(8) + ",\n" +
- " \"affine_to_cpu\" : false, \"strided_index\" : 1, \"mem_chunk\" : 2,\n"
+ " \"affine_to_cpu\" : false, \"strided_index\" : " + (ctx.isNVIDIA ? "0" : "1") + ", \"mem_chunk\" : 2,\n"
" \"comp_mode\" : true\n" +
" },\n";
}
@@ -148,11 +166,13 @@ private:
configTpl.replace("PLATFORMINDEX",std::to_string(platformIndex));
configTpl.replace("GPUCONFIG",conf);
configTpl.write(params::inst().configFileAMD);
- printer::inst()->print_msg(L0, "AMD: GPU configuration stored in file '%s'", params::inst().configFileAMD.c_str());
+
+ const std::string backendName = xmrstak::params::inst().openCLVendor;
+ printer::inst()->print_msg(L0, "%s: GPU (OpenCL) configuration stored in file '%s'", backendName.c_str(), params::inst().configFileAMD.c_str());
}
std::vector<GpuContext> devVec;
};
} // namespace amd
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp
index cab5ad9dac159699565e4f6427def89631b2f204..4353e3d05d21aa01fb76be8139ff5ffe340432cd 100644
--- a/xmrstak/backend/amd/minethd.cpp
+++ b/xmrstak/backend/amd/minethd.cpp
@@ -137,6 +137,8 @@ std::vector<iBackend*>* minethd::thread_starter(uint32_t threadOffset, miner_wor
for (i = 0; i < n; i++)
{
jconf::inst()->GetThreadConfig(i, cfg);
+
+ const std::string backendName = xmrstak::params::inst().openCLVendor;
if(cfg.cpu_aff >= 0)
{
@@ -144,10 +146,10 @@ std::vector<iBackend*>* minethd::thread_starter(uint32_t threadOffset, miner_wor
printer::inst()->print_msg(L1, "WARNING on macOS thread affinity is only advisory.");
#endif
- printer::inst()->print_msg(L1, "Starting AMD GPU thread %d, affinity: %d.", i, (int)cfg.cpu_aff);
+ printer::inst()->print_msg(L1, "Starting %s GPU (OpenCL) thread %d, affinity: %d.", backendName.c_str(), i, (int)cfg.cpu_aff);
}
else
- printer::inst()->print_msg(L1, "Starting AMD GPU thread %d, no affinity.", i);
+ printer::inst()->print_msg(L1, "Starting %s GPU (OpenCL) thread %d, no affinity.", backendName.c_str(), i);
minethd* thd = new minethd(pWork, i + threadOffset, &vGpuData[i], cfg);
pvThreads->push_back(thd);
@@ -193,12 +195,13 @@ void minethd::work_main()
cpu_ctx = cpu::minethd::minethd_alloc_ctx();
// start with root algorithm and switch later if fork version is reached
- auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+ auto miner_algo = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot();
cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
globalStates::inst().iConsumeCnt++;
uint8_t version = 0;
+ size_t lastPoolId = 0;
while (bQuit == 0)
{
@@ -217,13 +220,20 @@ void minethd::work_main()
}
uint8_t new_version = oWork.getVersion();
- if(new_version != version)
+ if(new_version != version || oWork.iPoolId != lastPoolId)
{
- if(new_version >= ::jconf::inst()->GetMiningForkVersion())
+ coinDescription coinDesc = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(oWork.iPoolId);
+ if(new_version >= coinDesc.GetMiningForkVersion())
+ {
+ miner_algo = coinDesc.GetMiningAlgo();
+ hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
+ }
+ else
{
- miner_algo = ::jconf::inst()->GetMiningAlgo();
+ miner_algo = coinDesc.GetMiningAlgoRoot();
hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
}
+ lastPoolId = oWork.iPoolId;
version = new_version;
}
@@ -233,7 +243,7 @@ void minethd::work_main()
assert(sizeof(job_result::sJobID) == sizeof(pool_job::sJobID));
uint64_t target = oWork.iTarget;
- XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target, miner_algo, version);
+ XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target, miner_algo);
if(oWork.bNiceHash)
pGpuCtx->Nonce = *(uint32_t*)(oWork.bWorkBlob + 39);
@@ -249,7 +259,7 @@ void minethd::work_main()
cl_uint results[0x100];
memset(results,0,sizeof(cl_uint)*(0x100));
- XMRRunJob(pGpuCtx, results, miner_algo, version);
+ XMRRunJob(pGpuCtx, results, miner_algo);
for(size_t i = 0; i < results[0xFF]; i++)
{
diff --git a/xmrstak/backend/backendConnector.cpp b/xmrstak/backend/backendConnector.cpp
index d735cb391c2240572365cf1ad70ea83ee1342daa..6f80a0f73ad45fd818a23f55423b327a9e869966 100644
--- a/xmrstak/backend/backendConnector.cpp
+++ b/xmrstak/backend/backendConnector.cpp
@@ -77,11 +77,12 @@ std::vector<iBackend*>* BackendConnector::thread_starter(miner_work& pWork)
#ifndef CONF_NO_OPENCL
if(params::inst().useAMD)
{
- plugin amdplugin("AMD", "xmrstak_opencl_backend");
+ const std::string backendName = xmrstak::params::inst().openCLVendor;
+ plugin amdplugin(backendName, "xmrstak_opencl_backend");
std::vector<iBackend*>* amdThreads = amdplugin.startBackend(static_cast<uint32_t>(pvThreads->size()), pWork, environment::inst());
pvThreads->insert(std::end(*pvThreads), std::begin(*amdThreads), std::end(*amdThreads));
if(amdThreads->size() == 0)
- printer::inst()->print_msg(L0, "WARNING: backend AMD disabled.");
+ printer::inst()->print_msg(L0, "WARNING: backend %s (OpenCL) disabled.", backendName.c_str());
}
#endif
@@ -99,4 +100,4 @@ std::vector<iBackend*>* BackendConnector::thread_starter(miner_work& pWork)
return pvThreads;
}
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/backendConnector.hpp b/xmrstak/backend/backendConnector.hpp
index da3dc7782cded3ac57d53b4796b9bbe85bd2b2c6..66d873e4853da9fd203c927d47848be5ca95874a 100644
--- a/xmrstak/backend/backendConnector.hpp
+++ b/xmrstak/backend/backendConnector.hpp
@@ -18,4 +18,4 @@ namespace xmrstak
static bool self_test();
};
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/cpu/autoAdjust.hpp b/xmrstak/backend/cpu/autoAdjust.hpp
index 969d47825384533dfd823ba98160ce6f1860354d..518721a2f6d423176a97fc726e2cac13b5285922 100644
--- a/xmrstak/backend/cpu/autoAdjust.hpp
+++ b/xmrstak/backend/cpu/autoAdjust.hpp
@@ -36,7 +36,10 @@ public:
bool printConfig()
{
- const size_t hashMemSizeKB = cn_select_memory(::jconf::inst()->GetMiningAlgo()) / 1024u;
+ const size_t hashMemSizeKB = std::max(
+ cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()),
+ cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot())
+ ) / 1024u;
const size_t halfHashMemSizeKB = hashMemSizeKB / 2u;
configEditor configTpl{};
@@ -172,4 +175,4 @@ private:
};
} // namespace cpu
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/cpu/autoAdjustHwloc.hpp b/xmrstak/backend/cpu/autoAdjustHwloc.hpp
index 568abb5cd6a86ef7b12fb8cd406bbb561ddbaa47..b1f3914735092761bf18b7331359cae70df40193 100644
--- a/xmrstak/backend/cpu/autoAdjustHwloc.hpp
+++ b/xmrstak/backend/cpu/autoAdjustHwloc.hpp
@@ -28,7 +28,10 @@ public:
autoAdjust()
{
- hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
+ hashMemSize = std::max(
+ cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()),
+ cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot())
+ );
halfHashMemSize = hashMemSize / 2u;
}
@@ -214,4 +217,4 @@ private:
};
} // namespace cpu
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/cpu/crypto/c_blake256.c b/xmrstak/backend/cpu/crypto/c_blake256.c
index ff623ddb8769688f78bca97617b2c1aae2097875..e5fadfe74bf491d6469f32d7b948e908454899e8 100644
--- a/xmrstak/backend/cpu/crypto/c_blake256.c
+++ b/xmrstak/backend/cpu/crypto/c_blake256.c
@@ -124,7 +124,7 @@ void blake224_init(state *S) {
}
// datalen = number of bits
-void blake256_update(state *S, const uint8_t *data, uint64_t datalen) {
+void blake256_update(state *S, const uint8_t *data, uint32_t datalen) {
int left = S->buflen >> 3;
int fill = 64 - left;
@@ -155,7 +155,7 @@ void blake256_update(state *S, const uint8_t *data, uint64_t datalen) {
}
// datalen = number of bits
-void blake224_update(state *S, const uint8_t *data, uint64_t datalen) {
+void blake224_update(state *S, const uint8_t *data, uint32_t datalen) {
blake256_update(S, data, datalen);
}
@@ -206,7 +206,7 @@ void blake224_final(state *S, uint8_t *digest) {
}
// inlen = number of bytes
-void blake256_hash(uint8_t *out, const uint8_t *in, uint64_t inlen) {
+void blake256_hash(uint8_t *out, const uint8_t *in, uint32_t inlen) {
state S;
blake256_init(&S);
blake256_update(&S, in, inlen * 8);
@@ -214,7 +214,7 @@ void blake256_hash(uint8_t *out, const uint8_t *in, uint64_t inlen) {
}
// inlen = number of bytes
-void blake224_hash(uint8_t *out, const uint8_t *in, uint64_t inlen) {
+void blake224_hash(uint8_t *out, const uint8_t *in, uint32_t inlen) {
state S;
blake224_init(&S);
blake224_update(&S, in, inlen * 8);
@@ -282,13 +282,13 @@ void hmac_blake224_init(hmac_state *S, const uint8_t *_key, uint64_t keylen) {
}
// datalen = number of bits
-void hmac_blake256_update(hmac_state *S, const uint8_t *data, uint64_t datalen) {
+void hmac_blake256_update(hmac_state *S, const uint8_t *data, uint32_t datalen) {
// update the inner state
blake256_update(&S->inner, data, datalen);
}
// datalen = number of bits
-void hmac_blake224_update(hmac_state *S, const uint8_t *data, uint64_t datalen) {
+void hmac_blake224_update(hmac_state *S, const uint8_t *data, uint32_t datalen) {
// update the inner state
blake224_update(&S->inner, data, datalen);
}
@@ -310,7 +310,7 @@ void hmac_blake224_final(hmac_state *S, uint8_t *digest) {
}
// keylen = number of bytes; inlen = number of bytes
-void hmac_blake256_hash(uint8_t *out, const uint8_t *key, uint64_t keylen, const uint8_t *in, uint64_t inlen) {
+void hmac_blake256_hash(uint8_t *out, const uint8_t *key, uint64_t keylen, const uint8_t *in, uint32_t inlen) {
hmac_state S;
hmac_blake256_init(&S, key, keylen);
hmac_blake256_update(&S, in, inlen * 8);
@@ -318,7 +318,7 @@ void hmac_blake256_hash(uint8_t *out, const uint8_t *key, uint64_t keylen, const
}
// keylen = number of bytes; inlen = number of bytes
-void hmac_blake224_hash(uint8_t *out, const uint8_t *key, uint64_t keylen, const uint8_t *in, uint64_t inlen) {
+void hmac_blake224_hash(uint8_t *out, const uint8_t *key, uint64_t keylen, const uint8_t *in, uint32_t inlen) {
hmac_state S;
hmac_blake224_init(&S, key, keylen);
hmac_blake224_update(&S, in, inlen * 8);
diff --git a/xmrstak/backend/cpu/crypto/c_blake256.h b/xmrstak/backend/cpu/crypto/c_blake256.h
index b9c2aad0dfd71f9624467cc547dfae0038b7b705..06c7917af8c159c85d425a7dd7b3afc87f9ec9d7 100644
--- a/xmrstak/backend/cpu/crypto/c_blake256.h
+++ b/xmrstak/backend/cpu/crypto/c_blake256.h
@@ -17,27 +17,27 @@ typedef struct {
void blake256_init(state *);
void blake224_init(state *);
-void blake256_update(state *, const uint8_t *, uint64_t);
-void blake224_update(state *, const uint8_t *, uint64_t);
+void blake256_update(state *, const uint8_t *, uint32_t);
+void blake224_update(state *, const uint8_t *, uint32_t);
void blake256_final(state *, uint8_t *);
void blake224_final(state *, uint8_t *);
-void blake256_hash(uint8_t *, const uint8_t *, uint64_t);
-void blake224_hash(uint8_t *, const uint8_t *, uint64_t);
+void blake256_hash(uint8_t *, const uint8_t *, uint32_t);
+void blake224_hash(uint8_t *, const uint8_t *, uint32_t);
/* HMAC functions: */
void hmac_blake256_init(hmac_state *, const uint8_t *, uint64_t);
void hmac_blake224_init(hmac_state *, const uint8_t *, uint64_t);
-void hmac_blake256_update(hmac_state *, const uint8_t *, uint64_t);
-void hmac_blake224_update(hmac_state *, const uint8_t *, uint64_t);
+void hmac_blake256_update(hmac_state *, const uint8_t *, uint32_t);
+void hmac_blake224_update(hmac_state *, const uint8_t *, uint32_t);
void hmac_blake256_final(hmac_state *, uint8_t *);
void hmac_blake224_final(hmac_state *, uint8_t *);
-void hmac_blake256_hash(uint8_t *, const uint8_t *, uint64_t, const uint8_t *, uint64_t);
-void hmac_blake224_hash(uint8_t *, const uint8_t *, uint64_t, const uint8_t *, uint64_t);
+void hmac_blake256_hash(uint8_t *, const uint8_t *, uint64_t, const uint8_t *, uint32_t);
+void hmac_blake224_hash(uint8_t *, const uint8_t *, uint64_t, const uint8_t *, uint32_t);
#endif /* _BLAKE256_H_ */
diff --git a/xmrstak/backend/cpu/crypto/c_skein.h b/xmrstak/backend/cpu/crypto/c_skein.h
index 6165a2ace27d7f34c71f9d998255a00e9dec01d7..86dbc08020ede237b0dba1e122eb8c43ba6913af 100644
--- a/xmrstak/backend/cpu/crypto/c_skein.h
+++ b/xmrstak/backend/cpu/crypto/c_skein.h
@@ -37,7 +37,7 @@ typedef enum
}
SkeinHashReturn;
-typedef size_t SkeinDataLength; /* bit count type */
+typedef uint32_t SkeinDataLength; /* bit count type */
typedef u08b_t SkeinBitSequence; /* bit stream type */
/* "all-in-one" call */
diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
index 5203de8606c25dd7e472cde72417cc3fa9c40226..7562de1bff3cb98e932bea68725cfefe51ee9a31 100644
--- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
+++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
@@ -43,7 +43,7 @@ extern "C"
{
void keccak(const uint8_t *in, int inlen, uint8_t *md, int mdlen);
void keccakf(uint64_t st[25], int rounds);
- extern void(*const extra_hashes[4])(const void *, size_t, char *);
+ extern void(*const extra_hashes[4])(const void *, uint32_t, char *);
}
// This will shift and xor tmp1 into itself as 4 32-bit vals such as
@@ -429,7 +429,7 @@ inline void cryptonight_monero_tweak(uint64_t* mem_out, __m128i tmp)
tmp = _mm_castps_si128(_mm_movehl_ps(_mm_castsi128_ps(tmp), _mm_castsi128_ps(tmp)));
uint64_t vh = _mm_cvtsi128_si64(tmp);
- uint8_t x = vh >> 24;
+ uint8_t x = static_cast<uint8_t>(vh >> 24);
static const uint16_t table = 0x7531;
const uint8_t index = (((x >> 3) & 6) | (x & 1)) << 1;
vh ^= ((table >> index) & 0x3) << 28;
diff --git a/xmrstak/backend/cpu/crypto/cryptonight_common.cpp b/xmrstak/backend/cpu/crypto/cryptonight_common.cpp
index 17fa24b0648d9e0be72036847958f827b7ba97d9..ee3b66301e0a0f09d75f0d0e0efe504f4f0d62c2 100644
--- a/xmrstak/backend/cpu/crypto/cryptonight_common.cpp
+++ b/xmrstak/backend/cpu/crypto/cryptonight_common.cpp
@@ -35,6 +35,7 @@ extern "C"
#include "xmrstak/jconf.hpp"
#include <stdio.h>
#include <stdlib.h>
+#include <algorithm>
#ifdef __GNUC__
#include <mm_malloc.h>
@@ -55,23 +56,23 @@ extern "C"
#include <string.h>
#endif // _WIN32
-void do_blake_hash(const void* input, size_t len, char* output) {
+void do_blake_hash(const void* input, uint32_t len, char* output) {
blake256_hash((uint8_t*)output, (const uint8_t*)input, len);
}
-void do_groestl_hash(const void* input, size_t len, char* output) {
+void do_groestl_hash(const void* input, uint32_t len, char* output) {
groestl((const uint8_t*)input, len * 8, (uint8_t*)output);
}
-void do_jh_hash(const void* input, size_t len, char* output) {
+void do_jh_hash(const void* input, uint32_t len, char* output) {
jh_hash(32 * 8, (const uint8_t*)input, 8 * len, (uint8_t*)output);
}
-void do_skein_hash(const void* input, size_t len, char* output) {
+void do_skein_hash(const void* input, uint32_t len, char* output) {
skein_hash(8 * 32, (const uint8_t*)input, 8 * len, (uint8_t*)output);
}
-void (* const extra_hashes[4])(const void *, size_t, char *) = {do_blake_hash, do_groestl_hash, do_jh_hash, do_skein_hash};
+void (* const extra_hashes[4])(const void *, uint32_t, char *) = {do_blake_hash, do_groestl_hash, do_jh_hash, do_skein_hash};
#ifdef _WIN32
#include "xmrstak/misc/uac.hpp"
@@ -202,7 +203,10 @@ size_t cryptonight_init(size_t use_fast_mem, size_t use_mlock, alloc_msg* msg)
cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, alloc_msg* msg)
{
- size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
+ size_t hashMemSize = std::max(
+ cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()),
+ cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot())
+ );
cryptonight_ctx* ptr = (cryptonight_ctx*)_mm_malloc(sizeof(cryptonight_ctx), 4096);
@@ -278,7 +282,10 @@ cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, al
void cryptonight_free_ctx(cryptonight_ctx* ctx)
{
- size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
+ size_t hashMemSize = std::max(
+ cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()),
+ cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot())
+ );
if(ctx->ctx_info[0] != 0)
{
diff --git a/xmrstak/backend/cpu/crypto/hash.h b/xmrstak/backend/cpu/crypto/hash.h
index c12d355f49d68c03e829f57b6c7f7047c08c356c..2af330932eddc362e6fa0b648ff571f62b66355c 100644
--- a/xmrstak/backend/cpu/crypto/hash.h
+++ b/xmrstak/backend/cpu/crypto/hash.h
@@ -1,5 +1,7 @@
#pragma once
+#include <stdint.h>
+
typedef unsigned char BitSequence;
-typedef unsigned long long DataLength;
+typedef uint32_t DataLength;
typedef enum {SUCCESS = 0, FAIL = 1, BAD_HASHLEN = 2} HashReturn;
diff --git a/xmrstak/backend/cpu/crypto/skein_port.h b/xmrstak/backend/cpu/crypto/skein_port.h
index 9cbefcb1a79f1715a84cfb253a80478ca3238b4e..99641bcdf9227a2cd6c971a652c44e58c7f6e0e5 100644
--- a/xmrstak/backend/cpu/crypto/skein_port.h
+++ b/xmrstak/backend/cpu/crypto/skein_port.h
@@ -49,7 +49,7 @@
multiple of size / 8)
ptr_cast(x,size) casts a pointer to a pointer to a
- varaiable of length 'size' bits
+ variable of length 'size' bits
*/
#define ui_type(size) uint##size##_t
diff --git a/xmrstak/backend/cpu/hwlocMemory.cpp b/xmrstak/backend/cpu/hwlocMemory.cpp
index 94d2b53f78dba683a6a01623b1e3be48e2215e3b..089570fc01bfa137d3d7ae55093851f480a260f3 100644
--- a/xmrstak/backend/cpu/hwlocMemory.cpp
+++ b/xmrstak/backend/cpu/hwlocMemory.cpp
@@ -30,7 +30,7 @@ void bindMemoryToNUMANode( size_t puId )
depth = hwloc_get_type_depth(topology, HWLOC_OBJ_PU);
- for( size_t i = 0;
+ for( uint32_t i = 0;
i < hwloc_get_nbobjs_by_depth(topology, depth);
i++ )
{
diff --git a/xmrstak/backend/cpu/jconf.cpp b/xmrstak/backend/cpu/jconf.cpp
index 399dd164792f75185d0343202b2b2319f0d57c05..49da7ae2d2b575421bd227af5663b49f2799c924 100644
--- a/xmrstak/backend/cpu/jconf.cpp
+++ b/xmrstak/backend/cpu/jconf.cpp
@@ -259,4 +259,4 @@ bool jconf::parse_config(const char* sFilename)
}
} // namespace cpu
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/cpu/jconf.hpp b/xmrstak/backend/cpu/jconf.hpp
index e98ed160a99ec9647b3e5b9ad8eaae521d0d1423..be855036eceaeb32f1e7fb8094dd1b244fb29cc7 100644
--- a/xmrstak/backend/cpu/jconf.hpp
+++ b/xmrstak/backend/cpu/jconf.hpp
@@ -40,4 +40,4 @@ private:
};
} // namespace cpu
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp
index 5b56f851cf22175ab6343a161e330debb5c2978e..f8f70f9702b7b1e47d4418ad561ab1b1ef5bbd81 100644
--- a/xmrstak/backend/cpu/minethd.cpp
+++ b/xmrstak/backend/cpu/minethd.cpp
@@ -231,7 +231,7 @@ bool minethd::self_test()
bool bResult = true;
- if(::jconf::inst()->GetMiningAlgo() == cryptonight)
+ if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight)
{
unsigned char out[32 * MAX_N];
cn_hash_fun hashf;
@@ -276,13 +276,13 @@ bool minethd::self_test()
"\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05"
"\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 160) == 0;
}
- else if(::jconf::inst()->GetMiningAlgo() == cryptonight_lite)
+ else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_lite)
{
}
- else if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero)
+ else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_monero)
{
}
- else if(::jconf::inst()->GetMiningAlgo() == cryptonight_aeon)
+ else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_aeon)
{
}
@@ -424,7 +424,7 @@ void minethd::work_main()
job_result result;
// start with root algorithm and switch later if fork version is reached
- auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+ auto miner_algo = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot();
cn_hash_fun hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
ctx = minethd_alloc_ctx();
@@ -434,6 +434,7 @@ void minethd::work_main()
result.iThreadId = iThreadNo;
uint8_t version = 0;
+ size_t lastPoolId = 0;
while (bQuit == 0)
{
@@ -461,13 +462,20 @@ void minethd::work_main()
result.iNonce = *piNonce;
uint8_t new_version = oWork.getVersion();
- if(new_version != version)
+ if(new_version != version || oWork.iPoolId != lastPoolId)
{
- if(new_version >= ::jconf::inst()->GetMiningForkVersion())
+ coinDescription coinDesc = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(oWork.iPoolId);
+ if(new_version >= coinDesc.GetMiningForkVersion())
{
- miner_algo = ::jconf::inst()->GetMiningAlgo();
+ miner_algo = coinDesc.GetMiningAlgo();
hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
}
+ else
+ {
+ miner_algo = coinDesc.GetMiningAlgoRoot();
+ hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
+ }
+ lastPoolId = oWork.iPoolId;
version = new_version;
}
@@ -627,22 +635,22 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
void minethd::double_work_main()
{
- multiway_work_main<2>();
+ multiway_work_main<2u>();
}
void minethd::triple_work_main()
{
- multiway_work_main<3>();
+ multiway_work_main<3u>();
}
void minethd::quad_work_main()
{
- multiway_work_main<4>();
+ multiway_work_main<4u>();
}
void minethd::penta_work_main()
{
- multiway_work_main<5>();
+ multiway_work_main<5u>();
}
template<size_t N>
@@ -656,7 +664,7 @@ void minethd::prep_multiway_work(uint8_t *bWorkBlob, uint32_t **piNonce)
}
}
-template<size_t N>
+template<uint32_t N>
void minethd::multiway_work_main()
{
if(affinity >= 0) //-1 means no affinity
@@ -689,9 +697,10 @@ void minethd::multiway_work_main()
globalStates::inst().iConsumeCnt++;
// start with root algorithm and switch later if fork version is reached
- auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+ auto miner_algo = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot();
cn_hash_fun_multi hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
uint8_t version = 0;
+ size_t lastPoolId = 0;
while (bQuit == 0)
{
@@ -718,13 +727,20 @@ void minethd::multiway_work_main()
iNonce = *piNonce[0];
uint8_t new_version = oWork.getVersion();
- if(new_version != version)
+ if(new_version != version || oWork.iPoolId != lastPoolId)
{
- if(new_version >= ::jconf::inst()->GetMiningForkVersion())
+ coinDescription coinDesc = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(oWork.iPoolId);
+ if(new_version >= coinDesc.GetMiningForkVersion())
+ {
+ miner_algo = coinDesc.GetMiningAlgo();
+ hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
+ }
+ else
{
- miner_algo = ::jconf::inst()->GetMiningAlgo();
+ miner_algo = coinDesc.GetMiningAlgoRoot();
hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
}
+ lastPoolId = oWork.iPoolId;
version = new_version;
}
@@ -769,4 +785,4 @@ void minethd::multiway_work_main()
}
} // namespace cpu
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/cpu/minethd.hpp b/xmrstak/backend/cpu/minethd.hpp
index cd78343081e480f7f472d0b937fb9c6f9c44ca62..85a95d1598c13c4458bb10385a0bcea981589591 100644
--- a/xmrstak/backend/cpu/minethd.hpp
+++ b/xmrstak/backend/cpu/minethd.hpp
@@ -35,7 +35,7 @@ private:
minethd(miner_work& pWork, size_t iNo, int iMultiway, bool no_prefetch, int64_t affinity);
- template<size_t N>
+ template<uint32_t N>
void multiway_work_main();
template<size_t N>
@@ -65,4 +65,4 @@ private:
};
} // namespace cpu
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/globalStates.cpp b/xmrstak/backend/globalStates.cpp
index 78823c53cab142f92ffa46ba9e306530b0ee10b2..1ec79839d750ef173b923bb92e61653a44dae54e 100644
--- a/xmrstak/backend/globalStates.cpp
+++ b/xmrstak/backend/globalStates.cpp
@@ -53,4 +53,4 @@ void globalStates::switch_work(miner_work& pWork, pool_data& dat)
iGlobalJobNo++;
}
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/globalStates.hpp b/xmrstak/backend/globalStates.hpp
index 442be699c755298f75b7c15b87dfa188b0721c3c..fafd232d95fe0fe3f36bde17cb3889b330834c80 100644
--- a/xmrstak/backend/globalStates.hpp
+++ b/xmrstak/backend/globalStates.hpp
@@ -55,4 +55,4 @@ private:
}
};
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/iBackend.hpp b/xmrstak/backend/iBackend.hpp
index 3d2115a19a7256b2d2deee85ec519f32601b6ca1..fdc647e08d30e9c4ae0dc9eac19458e58ed4b853 100644
--- a/xmrstak/backend/iBackend.hpp
+++ b/xmrstak/backend/iBackend.hpp
@@ -47,4 +47,4 @@ namespace xmrstak
}
};
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/miner_work.hpp b/xmrstak/backend/miner_work.hpp
index 9e5a4e46408120f2828b09e25706af8ae8ed9dfa..438ec0d900472d7b60a0bc5a19c378dfd9d137a7 100644
--- a/xmrstak/backend/miner_work.hpp
+++ b/xmrstak/backend/miner_work.hpp
@@ -81,4 +81,4 @@ namespace xmrstak
}
};
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/nvidia/autoAdjust.hpp b/xmrstak/backend/nvidia/autoAdjust.hpp
index be7d1ce433c65824bf4b05d503e2c80dbdade15b..d8bb6215fec26c0c06e6625b4d002b9af2445764 100644
--- a/xmrstak/backend/nvidia/autoAdjust.hpp
+++ b/xmrstak/backend/nvidia/autoAdjust.hpp
@@ -50,7 +50,7 @@ public:
ctx.device_blocks = -1;
ctx.device_threads = -1;
- // set all evice option those marked as auto (-1) to a valid value
+ // set all device option those marked as auto (-1) to a valid value
#ifndef _WIN32
ctx.device_bfactor = 0;
ctx.device_bsleep = 0;
@@ -109,4 +109,4 @@ private:
};
} // namespace nvidia
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/nvidia/jconf.hpp b/xmrstak/backend/nvidia/jconf.hpp
index 7f60f1d907c34402f6552d439b0f894274c635cc..b4ebaa03549c003ae82b1d7e96e524070f8d47d6 100644
--- a/xmrstak/backend/nvidia/jconf.hpp
+++ b/xmrstak/backend/nvidia/jconf.hpp
@@ -49,4 +49,4 @@ private:
};
} // namespace nvidia
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp
index 804c06a711ee8657b0d701459ca23cad9039f23d..92f5f789b0c649f4aa2a1270a8612a535a3b47ad 100644
--- a/xmrstak/backend/nvidia/minethd.cpp
+++ b/xmrstak/backend/nvidia/minethd.cpp
@@ -239,7 +239,7 @@ void minethd::work_main()
cpu_ctx = cpu::minethd::minethd_alloc_ctx();
// start with root algorithm and switch later if fork version is reached
- auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+ auto miner_algo = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot();
cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
uint32_t iNonce;
@@ -247,6 +247,7 @@ void minethd::work_main()
globalStates::inst().iConsumeCnt++;
uint8_t version = 0;
+ size_t lastPoolId = 0;
while (bQuit == 0)
{
@@ -264,13 +265,20 @@ void minethd::work_main()
continue;
}
uint8_t new_version = oWork.getVersion();
- if(new_version != version)
+ if(new_version != version || oWork.iPoolId != lastPoolId)
{
- if(new_version >= ::jconf::inst()->GetMiningForkVersion())
+ coinDescription coinDesc = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(oWork.iPoolId);
+ if(new_version >= coinDesc.GetMiningForkVersion())
{
- miner_algo = ::jconf::inst()->GetMiningAlgo();
+ miner_algo = coinDesc.GetMiningAlgo();
hash_fun = cpu::minethd::func_selector(::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);
+ }
+ lastPoolId = oWork.iPoolId;
version = new_version;
}
diff --git a/xmrstak/backend/nvidia/minethd.hpp b/xmrstak/backend/nvidia/minethd.hpp
index fcd24fa89e12ba9f70ce238dc11f057f446a5912..89c294444cdf369acfc35247a5ee60618dff3e80 100644
--- a/xmrstak/backend/nvidia/minethd.hpp
+++ b/xmrstak/backend/nvidia/minethd.hpp
@@ -60,4 +60,4 @@ private:
};
} // namespace nvidia
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
index 02c157ed70d8c212edc9a20e346620fb4a47b135..f192f01ddc4e0ced551ff72e3fa93e7bdfebf941 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
@@ -279,12 +279,15 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx)
if(gpuArch < 70)
CUDA_CHECK(ctx->device_id, cudaDeviceSetCacheConfig(cudaFuncCachePreferL1));
- size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
+ size_t hashMemSize = std::max(
+ cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()),
+ cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot())
+ );
size_t wsize = ctx->device_blocks * ctx->device_threads;
CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_state, 50 * sizeof(uint32_t) * wsize));
size_t ctx_b_size = 4 * sizeof(uint32_t) * wsize;
- if(cryptonight_heavy == ::jconf::inst()->GetMiningAlgo())
+ if(cryptonight_heavy == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo())
{
// extent ctx_b to hold the state of idx0
ctx_b_size += sizeof(uint32_t) * wsize;
@@ -458,7 +461,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
#undef XMRSTAK_PP_TOSTRING1
std::stringstream ss(archStringList);
- //transform string list sperated with `+` into a vector of integers
+ //transform string list separated with `+` into a vector of integers
int tmpArch;
while ( ss >> tmpArch )
arch.push_back( tmpArch );
@@ -492,7 +495,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
}
}
- // set all evice option those marked as auto (-1) to a valid value
+ // set all device option those marked as auto (-1) to a valid value
if(ctx->device_blocks == -1)
{
/* good values based of my experience
@@ -576,7 +579,10 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
ctx->total_device_memory = totalMemory;
ctx->free_device_memory = freeMemory;
- size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
+ size_t hashMemSize = std::max(
+ cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()),
+ cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot())
+ );
#ifdef WIN32
/* We use in windows bfactor (split slow kernel into smaller parts) to avoid
@@ -606,7 +612,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
// up to 16kibyte extra memory is used per thread for some kernel (lmem/local memory)
// 680bytes are extra meta data memory per hash
size_t perThread = hashMemSize + 16192u + 680u;
- if(cryptonight_heavy == ::jconf::inst()->GetMiningAlgo())
+ if(cryptonight_heavy == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo())
perThread += 50 * 4; // state double buffer
size_t max_intensity = limitedMemory / perThread;
diff --git a/xmrstak/backend/plugin.hpp b/xmrstak/backend/plugin.hpp
index 89cdf97f8154bf4d02cb2366e25b0b8c50cfb3c0..2610db8674b777d08cf62f69e3ddd7b75f771f83 100644
--- a/xmrstak/backend/plugin.hpp
+++ b/xmrstak/backend/plugin.hpp
@@ -109,4 +109,4 @@ struct plugin
* */
};
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/cli/cli-miner.cpp b/xmrstak/cli/cli-miner.cpp
index c425f04dcf3fa943210a36d6e0a3291c0d7bb972..6118682d8cde0fd00afd2de84492f67d9c10093f 100644
--- a/xmrstak/cli/cli-miner.cpp
+++ b/xmrstak/cli/cli-miner.cpp
@@ -80,6 +80,8 @@ void help()
#ifndef CONF_NO_OPENCL
cout<<" --noAMD disable the AMD miner backend"<<endl;
cout<<" --noAMDCache disable the AMD(OpenCL) cache for precompiled binaries"<<endl;
+ cout<<" --openCLVendor VENDOR use OpenCL driver of VENDOR and devices [AMD,NVIDIA]"<<endl;
+ cout<<" default: AMD"<<endl;
cout<<" --amd FILE AMD backend miner config file"<<endl;
#endif
#ifndef CONF_NO_CUDA
@@ -450,6 +452,24 @@ int main(int argc, char *argv[])
{
params::inst().useAMD = false;
}
+ else if(opName.compare("--openCLVendor") == 0)
+ {
+ ++i;
+ if( i >=argc )
+ {
+ printer::inst()->print_msg(L0, "No argument for parameter '--openCLVendor' given");
+ win_exit();
+ return 1;
+ }
+ std::string vendor(argv[i]);
+ params::inst().openCLVendor = vendor;
+ if(vendor != "AMD" && vendor != "NVIDIA")
+ {
+ printer::inst()->print_msg(L0, "'--openCLVendor' must be 'AMD' or 'NVIDIA'");
+ win_exit();
+ return 1;
+ }
+ }
else if(opName.compare("--noAMDCache") == 0)
{
params::inst().AMDCache = false;
diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp
index f99698a0e1a5b4a8945735431ed5c5bc3f9a4113..3af4028add0196935de7eee5e114db27d72dec22 100644
--- a/xmrstak/jconf.cpp
+++ b/xmrstak/jconf.cpp
@@ -86,30 +86,28 @@ configVal oConfigValues[] = {
constexpr size_t iConfigCnt = (sizeof(oConfigValues)/sizeof(oConfigValues[0]));
-struct xmrstak_coin_algo
-{
- const char* coin_name;
- xmrstak_algo algo;
- xmrstak_algo algo_root;
- uint8_t fork_version;
- const char* default_pool;
-};
-
-xmrstak_coin_algo coin_algos[] = {
- { "aeon7", cryptonight_aeon, cryptonight_lite, 7u, "mine.aeon-pool.com:5555" },
- { "cryptonight", cryptonight, cryptonight, 0u, nullptr },
- { "cryptonight_lite", cryptonight_lite, cryptonight_lite, 0u, nullptr },
- { "edollar", cryptonight, cryptonight, 0u, nullptr },
- { "electroneum", cryptonight, cryptonight, 0u, nullptr },
- { "graft", cryptonight, cryptonight, 0u, nullptr },
- { "intense", cryptonight, cryptonight, 0u, nullptr },
- { "karbo", cryptonight, cryptonight, 0u, nullptr },
- { "monero7", cryptonight_monero, cryptonight, 7u, "pool.usxmrpool.com:3333" },
- { "stellite", cryptonight_monero, cryptonight, 3u, nullptr },
- { "sumokoin", cryptonight_heavy, cryptonight, 3u, nullptr }
+xmrstak::coin_selection coins[] = {
+ // name, userpool, devpool, default_pool_suggestion
+ { "aeon7", {cryptonight_aeon, cryptonight_lite, 7u}, {cryptonight_aeon, cryptonight_lite, 7u}, "mine.aeon-pool.com:5555" },
+ { "bbscoin", {cryptonight_monero, cryptonight, 3u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+ { "croat", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+ { "cryptonight", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+ { "cryptonight_heavy", {cryptonight_heavy, cryptonight_heavy, 0u}, {cryptonight_heavy, cryptonight_heavy, 0u}, nullptr },
+ { "cryptonight_lite", {cryptonight_aeon, cryptonight_lite, 255u}, {cryptonight_aeon, cryptonight_lite, 7u}, nullptr },
+ { "cryptonight_lite_v7", {cryptonight_lite, cryptonight_aeon, 255u}, {cryptonight_aeon, cryptonight_lite, 7u}, nullptr },
+ { "cryptonight_v7", {cryptonight_monero, cryptonight_monero, 0u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+ { "edollar", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+ { "electroneum", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+ { "graft", {cryptonight_monero, cryptonight, 8u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+ { "haven", {cryptonight_heavy, cryptonight, 2u}, {cryptonight_heavy, cryptonight_heavy, 0u}, nullptr },
+ { "intense", {cryptonight_monero, cryptonight, 4u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+ { "karbo", {cryptonight_monero, cryptonight, 255u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+ { "monero7", {cryptonight_monero, cryptonight_monero, 0u}, {cryptonight_monero, cryptonight_monero, 0u}, "pool.usxmrpool.com:3333" },
+ { "stellite", {cryptonight_monero, cryptonight, 3u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+ { "sumokoin", {cryptonight_heavy, cryptonight_heavy, 0u}, {cryptonight_heavy, cryptonight_heavy, 0u}, nullptr }
};
-constexpr size_t coin_alogo_size = (sizeof(coin_algos)/sizeof(coin_algos[0]));
+constexpr size_t coin_alogo_size = (sizeof(coins)/sizeof(coins[0]));
inline bool checkType(Type have, Type want)
{
@@ -313,7 +311,7 @@ void jconf::GetAlgoList(std::string& list)
for(size_t i=0; i < coin_alogo_size; i++)
{
list += "\t- ";
- list += coin_algos[i].coin_name;
+ list += coins[i].coin_name;
list += "\n";
}
}
@@ -331,7 +329,7 @@ bool jconf::IsOnAlgoList(std::string& needle)
for(size_t i=0; i < coin_alogo_size; i++)
{
- if(needle == coin_algos[i].coin_name)
+ if(needle == coins[i].coin_name)
return true;
}
return false;
@@ -343,10 +341,10 @@ const char* jconf::GetDefaultPool(const char* needle)
for(size_t i=0; i < coin_alogo_size; i++)
{
- if(strcmp(needle, coin_algos[i].coin_name) == 0)
+ if(strcmp(needle, coins[i].coin_name) == 0)
{
- if(coin_algos[i].default_pool != nullptr)
- return coin_algos[i].default_pool;
+ if(coins[i].default_pool != nullptr)
+ return coins[i].default_pool;
else
return default_example;
}
@@ -630,16 +628,14 @@ bool jconf::parse_config(const char* sFilename, const char* sFilenamePools)
return false;
}
- if(ctmp == coin_algos[i].coin_name)
+ if(ctmp == coins[i].coin_name)
{
- mining_algo = coin_algos[i].algo;
- mining_algo_root = coin_algos[i].algo_root;
- mining_fork_version = coin_algos[i].fork_version;
+ currentCoin = coins[i];
break;
}
}
- if(mining_algo == invalid_algo)
+ if(currentCoin.GetDescription(1).GetMiningAlgo() == invalid_algo)
{
std::string cl;
GetAlgoList(cl);
diff --git a/xmrstak/jconf.hpp b/xmrstak/jconf.hpp
index 76b93ddaede4145dcac240708a95132310270670..102b70f54f761ca6f0add0c4b24b8430d69d53ef 100644
--- a/xmrstak/jconf.hpp
+++ b/xmrstak/jconf.hpp
@@ -1,7 +1,7 @@
#pragma once
-#include "xmrstak/backend/cryptonight.hpp"
#include "xmrstak/misc/environment.hpp"
+#include "xmrstak/misc/coinDescription.hpp"
#include "params.hpp"
#include <stdlib.h>
@@ -48,12 +48,8 @@ public:
bool TlsSecureAlgos();
- inline xmrstak_algo GetMiningAlgo() const { return mining_algo; }
+ inline xmrstak::coin_selection GetCurrentCoinSelection() const { return currentCoin; }
- inline xmrstak_algo GetMiningAlgoRoot() const { return mining_algo_root; }
-
- inline uint8_t GetMiningForkVersion() const { return mining_fork_version; }
-
std::string GetMiningCoin();
static void GetAlgoList(std::string& list);
@@ -94,7 +90,5 @@ private:
opaque_private* prv;
bool bHaveAes;
- xmrstak_algo mining_algo;
- xmrstak_algo mining_algo_root;
- uint8_t mining_fork_version;
+ xmrstak::coin_selection currentCoin;
};
diff --git a/xmrstak/misc/coinDescription.hpp b/xmrstak/misc/coinDescription.hpp
new file mode 100644
index 0000000000000000000000000000000000000000..73d9a95482bdb71b9707810c6f89834d17c12fa0
--- /dev/null
+++ b/xmrstak/misc/coinDescription.hpp
@@ -0,0 +1,60 @@
+#pragma once
+
+#include "xmrstak/backend/cryptonight.hpp"
+
+#include <stdlib.h>
+#include <string>
+
+
+namespace xmrstak
+{
+ struct coinDescription
+ {
+ xmrstak_algo algo;
+ xmrstak_algo algo_root;
+ uint8_t fork_version;
+
+ coinDescription() = default;
+
+ coinDescription(const xmrstak_algo in_algo, xmrstak_algo in_algo_root, const uint8_t in_fork_version) :
+ algo(in_algo), algo_root(in_algo_root), fork_version(in_fork_version)
+ {}
+
+ inline xmrstak_algo GetMiningAlgo() const { return algo; }
+ inline xmrstak_algo GetMiningAlgoRoot() const { return algo_root; }
+ inline uint8_t GetMiningForkVersion() const { return fork_version; }
+ };
+
+ struct coin_selection
+ {
+ const char* coin_name = nullptr;
+ /* [0] -> user pool
+ * [1] -> dev pool
+ */
+ coinDescription pool_coin[2];
+ const char* default_pool = nullptr;
+
+ coin_selection() = default;
+
+ coin_selection(
+ const char* in_coin_name,
+ const coinDescription user_coinDescription,
+ const coinDescription dev_coinDescription,
+ const char* in_default_pool
+ ) :
+ coin_name(in_coin_name), default_pool(in_default_pool)
+ {
+ pool_coin[0] = user_coinDescription;
+ pool_coin[1] = dev_coinDescription;
+ }
+
+ /** get coin description for the pool
+ *
+ * @param poolId 0 select dev pool, else the user pool is selected
+ */
+ inline coinDescription GetDescription(size_t poolId) const {
+ coinDescription tmp = (poolId == 0 ? pool_coin[1] : pool_coin[0]);
+ return tmp;
+ }
+ };
+} // namespace xmrstak
diff --git a/xmrstak/misc/configEditor.hpp b/xmrstak/misc/configEditor.hpp
index a840bc4f2462268cc8265c1e1f6c944134d8b6ad..8a81ad66db4115a056ae707d2f3f7f0b9d41c20a 100644
--- a/xmrstak/misc/configEditor.hpp
+++ b/xmrstak/misc/configEditor.hpp
@@ -54,4 +54,4 @@ struct configEditor
};
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/misc/console.cpp b/xmrstak/misc/console.cpp
index de5eed3db67f2fdc9c570ece18eba1db8c900526..d961b713ba51aa1da011cd29ded0dca0513dd388 100644
--- a/xmrstak/misc/console.cpp
+++ b/xmrstak/misc/console.cpp
@@ -225,7 +225,7 @@ void printer::print_str(const char* str)
//Do a press any key for the windows folk. *insert any key joke here*
#ifdef _WIN32
-void win_exit(size_t code)
+void win_exit(int code)
{
size_t envSize = 0;
getenv_s(&envSize, nullptr, 0, "XMRSTAK_NOWAIT");
@@ -238,7 +238,7 @@ void win_exit(size_t code)
}
#else
-void win_exit(size_t code)
+void win_exit(int code)
{
std::exit(code);
}
diff --git a/xmrstak/misc/console.hpp b/xmrstak/misc/console.hpp
index cfbeddd01ec806bd4dca8a441f58fc89e8f1883e..671763105fe70a97c9741d181d4077eabd8df455 100644
--- a/xmrstak/misc/console.hpp
+++ b/xmrstak/misc/console.hpp
@@ -49,4 +49,4 @@ private:
FILE* logfile;
};
-void win_exit(size_t code = 1);
+void win_exit(int code = 1);
diff --git a/xmrstak/misc/environment.hpp b/xmrstak/misc/environment.hpp
index 99c2db82c0f2d5c4542af28c86ff63032bd1a630..b67c858741171650ecd725938f3e31406e605c97 100644
--- a/xmrstak/misc/environment.hpp
+++ b/xmrstak/misc/environment.hpp
@@ -38,4 +38,4 @@ struct environment
params* pParams = nullptr;
};
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp
index 0e1dd9f4e203f8fe60f79e699d4b2d48988dc2ef..8c454ebc141af58f450c31c860795e206727b599 100644
--- a/xmrstak/misc/executor.cpp
+++ b/xmrstak/misc/executor.cpp
@@ -421,7 +421,9 @@ void executor::on_miner_result(size_t pool_id, job_result& oResult)
{
//Ignore errors silently
if(pool->is_running() && pool->is_logged_in())
- pool->cmd_submit(oResult.sJobID, oResult.iNonce, oResult.bResult, backend_name, backend_hashcount, total_hashcount, jconf::inst()->GetMiningAlgo());
+ pool->cmd_submit(oResult.sJobID, oResult.iNonce, oResult.bResult, backend_name,
+ backend_hashcount, total_hashcount, jconf::inst()->GetCurrentCoinSelection().GetDescription(0).GetMiningAlgo()
+ );
return;
}
@@ -432,7 +434,9 @@ void executor::on_miner_result(size_t pool_id, job_result& oResult)
}
size_t t_start = get_timestamp_ms();
- bool bResult = pool->cmd_submit(oResult.sJobID, oResult.iNonce, oResult.bResult, backend_name, backend_hashcount, total_hashcount, jconf::inst()->GetMiningAlgo());
+ bool bResult = pool->cmd_submit(oResult.sJobID, oResult.iNonce, oResult.bResult,
+ backend_name, backend_hashcount, total_hashcount, jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()
+ );
size_t t_len = get_timestamp_ms() - t_start;
if(t_len > 0xFFFF)
@@ -548,7 +552,7 @@ void executor::ex_main()
pools.emplace_back(i+1, params.poolURL.c_str(), params.poolUsername.c_str(), params.poolRigid.c_str(), params.poolPasswd.c_str(), 9.9, false, params.poolUseTls, "", params.nicehashMode);
}
- switch(jconf::inst()->GetMiningAlgo())
+ switch(jconf::inst()->GetCurrentCoinSelection().GetDescription(0).GetMiningAlgo())
{
case cryptonight_heavy:
if(dev_tls)
diff --git a/xmrstak/misc/executor.hpp b/xmrstak/misc/executor.hpp
index fbaa265173141c17764f11d7d29c8ad8db81906d..be5ee6c2f7b81f422e0c122a7b6a09b188b4530f 100644
--- a/xmrstak/misc/executor.hpp
+++ b/xmrstak/misc/executor.hpp
@@ -23,7 +23,7 @@ namespace cpu
class minethd;
} // namespace cpu
-} // namepsace xmrstak
+} // namespace xmrstak
class executor
{
@@ -54,7 +54,7 @@ private:
inline void set_timestamp() { dev_timestamp = get_timestamp(); };
- // In miliseconds, has to divide a second (1000ms) into an integer number
+ // In milliseconds, has to divide a second (1000ms) into an integer number
constexpr static size_t iTickTime = 500;
// Dev donation time period in seconds. 100 minutes by default.
@@ -64,7 +64,7 @@ private:
inline bool is_dev_time()
{
//Add 2 seconds to compensate for connect
- constexpr size_t dev_portion = double(iDevDonatePeriod) * fDevDonationLevel + 2;
+ constexpr size_t dev_portion = static_cast<size_t>(double(iDevDonatePeriod) * fDevDonationLevel + 2.);
if(dev_portion < 12) //No point in bothering with less than 10s
return false;
diff --git a/xmrstak/misc/telemetry.cpp b/xmrstak/misc/telemetry.cpp
index 738d287635a5879b753a6989b53c399a587145f9..a3a2a122fd5e31cc0af47d905c74e7dd64864fff 100644
--- a/xmrstak/misc/telemetry.cpp
+++ b/xmrstak/misc/telemetry.cpp
@@ -89,8 +89,8 @@ double telemetry::calc_telemetry_data(size_t iLastMilisec, size_t iThread)
return nan("");
double fHashes, fTime;
- fHashes = iLastestHashCnt - iEarliestHashCnt;
- fTime = iLastestStamp - iEarliestStamp;
+ fHashes = static_cast<double>(iLastestHashCnt - iEarliestHashCnt);
+ fTime = static_cast<double>(iLastestStamp - iEarliestStamp);
fTime /= 1000.0;
return fHashes / fTime;
@@ -105,4 +105,4 @@ void telemetry::push_perf_value(size_t iThd, uint64_t iHashCount, uint64_t iTime
iBucketTop[iThd] = (iTop + 1) & iBucketMask;
}
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/misc/telemetry.hpp b/xmrstak/misc/telemetry.hpp
index b35bbbfd6fc271fd7a4db5bef5bb91c78ecb1d1b..2f84dfa86f7bae77dc8d38ee354a9020708f2a01 100644
--- a/xmrstak/misc/telemetry.hpp
+++ b/xmrstak/misc/telemetry.hpp
@@ -21,4 +21,4 @@ private:
uint64_t** ppTimestamps;
};
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/misc/utility.cpp b/xmrstak/misc/utility.cpp
index 3b1369a36c4d870559da10f847355e27f1b1ee96..5177d14c23121f13e4069e4ace67acf6b1a39784 100644
--- a/xmrstak/misc/utility.cpp
+++ b/xmrstak/misc/utility.cpp
@@ -18,4 +18,4 @@ namespace xmrstak
}
);
}
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/misc/utility.hpp b/xmrstak/misc/utility.hpp
index b2e841dffae793b91441e14b8b1b2ee6821c7e83..8f2e99fb8819ac06a872d3c7591e524effa5ab6a 100644
--- a/xmrstak/misc/utility.hpp
+++ b/xmrstak/misc/utility.hpp
@@ -9,4 +9,4 @@ namespace xmrstak
* @return true if both strings are equal, else false
*/
bool strcmp_i(const std::string& str1, const std::string& str2);
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp
index d71aeb1a65976762b9fcc44efc27582f2ecf79f9..74d1c261f2f9e636470822381b0be722eabb1591 100644
--- a/xmrstak/net/jpsock.cpp
+++ b/xmrstak/net/jpsock.cpp
@@ -229,7 +229,7 @@ void jpsock::jpsock_thread()
else
disconnect_time = 0;
- std::unique_lock<std::mutex>(job_mutex);
+ std::unique_lock<std::mutex> lck(job_mutex);
memset(&oCurrentJob, 0, sizeof(oCurrentJob));
bRunning = false;
}
@@ -439,7 +439,7 @@ bool jpsock::process_pool_job(const opq_json_val* params)
if(motd != nullptr && motd->IsString() && (motd->GetStringLength() & 0x01) == 0)
{
- std::unique_lock<std::mutex>(motd_mutex);
+ std::unique_lock<std::mutex> lck(motd_mutex);
if(motd->GetStringLength() > 0)
{
pool_motd.resize(motd->GetStringLength()/2 + 1);
@@ -454,7 +454,7 @@ bool jpsock::process_pool_job(const opq_json_val* params)
executor::inst()->push_event(ex_event(oPoolJob, pool_id));
- std::unique_lock<std::mutex>(job_mutex);
+ std::unique_lock<std::mutex> lck(job_mutex);
oCurrentJob = oPoolJob;
return true;
}
@@ -679,13 +679,13 @@ bool jpsock::cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bRes
void jpsock::save_nonce(uint32_t nonce)
{
- std::unique_lock<std::mutex>(job_mutex);
+ std::unique_lock<std::mutex> lck(job_mutex);
oCurrentJob.iSavedNonce = nonce;
}
bool jpsock::get_current_job(pool_job& job)
{
- std::unique_lock<std::mutex>(job_mutex);
+ std::unique_lock<std::mutex> lck(job_mutex);
if(oCurrentJob.iWorkLen == 0)
return false;
@@ -699,7 +699,7 @@ bool jpsock::get_pool_motd(std::string& strin)
if(!ext_motd)
return false;
- std::unique_lock<std::mutex>(motd_mutex);
+ std::unique_lock<std::mutex> lck(motd_mutex);
if(pool_motd.size() > 0)
{
strin.assign(pool_motd);
diff --git a/xmrstak/net/msgstruct.hpp b/xmrstak/net/msgstruct.hpp
index 8c4bdbe11f25c05744e156b3789f560541e85899..e401f591716d84f9b16209ae3c40b1b77763ebe7 100644
--- a/xmrstak/net/msgstruct.hpp
+++ b/xmrstak/net/msgstruct.hpp
@@ -179,7 +179,7 @@ inline size_t get_timestamp()
return time_point_cast<seconds>(steady_clock::now()).time_since_epoch().count();
};
-//Get milisecond timestamp
+//Get millisecond timestamp
inline size_t get_timestamp_ms()
{
using namespace std::chrono;
diff --git a/xmrstak/net/socket.cpp b/xmrstak/net/socket.cpp
index 7c58a8e0674cedb5d9c9030cdfc337dda9878d8c..9bc608fa7d4a286124f3e0c2a4e3c3cac9724a83 100644
--- a/xmrstak/net/socket.cpp
+++ b/xmrstak/net/socket.cpp
@@ -156,7 +156,8 @@ int plain_socket::recv(char* buf, unsigned int len)
bool plain_socket::send(const char* buf)
{
- int pos = 0, slen = strlen(buf);
+ size_t pos = 0;
+ size_t slen = strlen(buf);
while (pos != slen)
{
int ret = ::send(hSocket, buf + pos, slen - pos, 0);
diff --git a/xmrstak/params.hpp b/xmrstak/params.hpp
index 3d584b95cae11d56c7202595f98a9f045c83251e..6e676cf411777ea584cc6fa1ec769bc1f8e705d0 100644
--- a/xmrstak/params.hpp
+++ b/xmrstak/params.hpp
@@ -24,6 +24,8 @@ struct params
bool AMDCache;
bool useNVIDIA;
bool useCPU;
+ // user selected OpenCL vendor
+ std::string openCLVendor;
bool poolUseTls = false;
std::string poolURL;
@@ -60,6 +62,7 @@ struct params
AMDCache(true),
useNVIDIA(true),
useCPU(true),
+ openCLVendor("AMD"),
configFile("config.txt"),
configFilePools("pools.txt"),
configFileAMD("amd.txt"),
@@ -69,4 +72,4 @@ struct params
};
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/pools.tpl b/xmrstak/pools.tpl
index f5afff6c5334798235e0eb37b46961eb1348d1f6..58ca8b5d630519d44c5b2939690d66cd5a06ed3d 100644
--- a/xmrstak/pools.tpl
+++ b/xmrstak/pools.tpl
@@ -21,16 +21,27 @@ POOLCONF],
* Currency to mine. Supported values:
*
* aeon7 (use this for Aeon's new PoW)
- * cryptonight (try this if your coin is not listed)
- * cryptonight_lite
+ * bbscoin (automatic switch with block version 3 to cryptonight_v7)
+ * croat
* edollar
* electroneum
* graft
+ * haven
* intense
* karbo
* monero7 (use this for Monero's new PoW)
- * sumokoin
+ * sumokoin (automatic switch with block version 3 to cryptonight_heavy)
+ *
+ * Native algorithms which not depends on any block versions:
*
+ * # 1MiB scratchpad memory
+ * cryptonight_lite
+ * cryptonight_lite_v7
+ * # 2MiB scratchpad memory
+ * cryptonight
+ * cryptonight_v7
+ * # 4MiB scratchpad memory
+ * cryptonight_heavy
*/
"currency" : "CURRENCY",
diff --git a/xmrstak/version.cpp b/xmrstak/version.cpp
index 579c27ac673b8ca98518dd20eeccff717d89f24b..f46060c8f7b168548cc0bc6bdd17b46705aa80a2 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.4.2"
+#define XMR_STAK_VERSION "2.4.3"
#if defined(_WIN32)
#define OS_TYPE "win"