From d676a22a34ee549587e105da7bb1c43e562e0a79 Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Mon, 3 Apr 2017 21:34:43 +0200 Subject: [PATCH 01/16] fix condition `global_worker_size` needs to be a multiple of `local_worker_size`, this is not always the case in the current code. Remove the bit magic with a easy to read `ceil`. --- amd_gpu/gpu.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/amd_gpu/gpu.c b/amd_gpu/gpu.c index 7571e5d..17c62db 100644 --- a/amd_gpu/gpu.c +++ b/amd_gpu/gpu.c @@ -675,7 +675,7 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput) return(ERR_OCL_API); } - BranchNonces[i] += BranchNonces[i] + (w_size - (BranchNonces[i] & (w_size - 1))); + BranchNonces[i] = ((size_t)ceil( (double)BranchNonces[i] / (double)w_size) ) * w_size; if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[i + 3], 1, &ctx->Nonce, BranchNonces + i, &w_size, 0, NULL, NULL)) != CL_SUCCESS) { printer_print_msg("Error %s when calling clEnqueueNDRangeKernel for kernel %d.", err_to_str(ret), i + 3); From a111e01218fa89ffffbe36a63139009bfe97bf1f Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Fri, 7 Apr 2017 10:01:43 +0200 Subject: [PATCH 02/16] add OpenCl 1.2 support guard unknown defined from OpenCl 2.X --- amd_gpu/gpu.c | 2 ++ 1 file changed, 2 insertions(+) diff --git a/amd_gpu/gpu.c b/amd_gpu/gpu.c index 7571e5d..57a20bf 100644 --- a/amd_gpu/gpu.c +++ b/amd_gpu/gpu.c @@ -163,10 +163,12 @@ const char* err_to_str(cl_int ret) return "CL_INVALID_LINKER_OPTIONS"; case CL_INVALID_DEVICE_PARTITION_COUNT: return "CL_INVALID_DEVICE_PARTITION_COUNT"; +#ifdef CL_VERSION_2_0 case CL_INVALID_PIPE_SIZE: return "CL_INVALID_PIPE_SIZE"; case CL_INVALID_DEVICE_QUEUE: return "CL_INVALID_DEVICE_QUEUE"; +#endif default: return "UNKNOWN_ERROR"; } From cd422ebe43275ee892b87a62d04002e501b9c46f Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Fri, 7 Apr 2017 11:23:06 +0200 Subject: [PATCH 03/16] refactor cmake - allow compiling without microhttpd - allow compiling without OpenSSL - use find opencl to provide platform independence - install `config.txt` with command `make install` - copy opencl folder on `make install` to the install directory --- CMakeLists.txt | 120 ++++++++++++++++++++++++++++++++++++++++++------- cli-miner.cpp | 7 ++- httpd.cpp | 4 ++ 3 files changed, 113 insertions(+), 18 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 67ab6ca..acf4093 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,38 +1,124 @@ project(xmr-stak-amd) -cmake_minimum_required(VERSION 2.8.10) +cmake_minimum_required(VERSION 3.1.3) +# enforce C++11 +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_CXX_EXTENSIONS OFF) +set(CMAKE_CXX_STANDARD 11) + +if(CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT) + set(CMAKE_INSTALL_PREFIX "${CMAKE_BINARY_DIR}" CACHE PATH "install prefix" FORCE) +endif(CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT) + +# help to find AMD app SDK on systems with a software module system +list(APPEND CMAKE_PREFIX_PATH "$ENV{AMDAPPSDKROOT}") +# allow user to extent CMAKE_PREFIX_PATH via environment variable +list(APPEND CMAKE_PREFIX_PATH "$ENV{CMAKE_PREFIX_PATH}") + +################################################################################ +# CMake user options +################################################################################ + +# gcc 5.1 is the first GNU version without CoW strings +# https://github.com/fireice-uk/xmr-stak-nvidia/pull/10#issuecomment-290821792 +# If you remove this guard to compile with older gcc versions the miner will produce +# a high rate of wrong shares. if("${CMAKE_CXX_COMPILER_ID}" STREQUAL "GNU") - if (CMAKE_CXX_COMPILER_VERSION VERSION_LESS 5.1) + if(CMAKE_CXX_COMPILER_VERSION VERSION_LESS 5.1) message(FATAL_ERROR "GCC version must be at least 5.1!") endif() endif() +set(BUILD_TYPE "Release;Debug") +if(NOT CMAKE_BUILD_TYPE) + set(CMAKE_BUILD_TYPE "Release" CACHE STRING "Choose the type of build" FORCE) +endif() +set_property(CACHE CMAKE_BUILD_TYPE PROPERTY STRINGS "${BUILD_TYPE}") + +############################################################################### +# Find OpenCL +############################################################################### + +find_package(OpenCL REQUIRED) +include_directories(SYSTEM ${OpenCL_INCLUDE_DIRS}) +set(LIBS ${LIBS} ${OpenCL_LIBRARY}) +link_directories(${OpenCL_LIBRARY}) + +################################################################################ +# Find PThreads +################################################################################ + +find_package(Threads REQUIRED) +set(LIBS ${LIBS} ${CMAKE_THREAD_LIBS_INIT}) + +################################################################################ +# Find microhttpd +################################################################################ + find_library(MHTD NAMES microhttpd) if("${MHTD}" STREQUAL "MHTD-NOTFOUND") - message(FATAL_ERROR "libmicrohttpd is required") + message(STATUS "microhttpd NOT found: disable http server") + add_definitions("-DCONF_NO_HTTPD") +else() + set(LIBS ${LIBS} ${MHTD}) endif() -find_package(OpenSSL REQUIRED) +############################################################################### +# Find OpenSSL +############################################################################### + +find_package(OpenSSL) include_directories(${OPENSSL_INCLUDE_DIR}) - -#set(CMAKE_VERBOSE_MAKEFILE ON) -set(CMAKE_CONFIGURATION_TYPES "RELEASE;STATIC") -if("${CMAKE_BUILD_TYPE}" STREQUAL "") - set(CMAKE_BUILD_TYPE RELEASE) +set(LIBS ${LIBS} ${OPENSSL_LIBRARIES}) +if(NOT OPENSSL_FOUND) + add_definitions("-DCONF_NO_TLS") endif() -set(CMAKE_C_FLAGS "-DNDEBUG -march=westmere -O3 -m64 -s") -set(CMAKE_CXX_FLAGS "${CMAKE_C_FLAGS} -std=c++11") +################################################################################ +# Compile & Link +################################################################################ -set(CMAKE_EXE_LINKER_FLAGS_RELSEASE "") -set(CMAKE_EXE_LINKER_FLAGS_STATIC "-static-libgcc -static-libstdc++") +include_directories(.) + +# activate sse2 and aes-ni +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -msse2 -maes") + +file(GLOB SRCFILES_CPP "*.cpp" "crypto/*.cpp") +file(GLOB SRCFILES_C "crypto/*.c" "amd_gpu/*.c") + +add_library(xmr-stak-c + STATIC + ${SRCFILES_C} +) +set_property(TARGET xmr-stak-c PROPERTY C_STANDARD 99) + +add_executable(xmr-stak-amd + ${SRCFILES_CPP} +) set(EXECUTABLE_OUTPUT_PATH "bin") +target_link_libraries(xmr-stak-amd xmr-stak-c ${LIBS}) -file(GLOB SOURCES "crypto/*.c" "crypto/*.cpp" "amd_gpu/*.c" "*.cpp") +################################################################################ +# Install +################################################################################ -add_executable(xmr-stak-amd ${SOURCES}) -target_link_libraries(xmr-stak-amd pthread microhttpd OpenCL crypto ssl) - +# 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-amd + RUNTIME DESTINATION "${CMAKE_INSTALL_PREFIX}/bin") +endif() + +install(DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/opencl" + DESTINATION "${CMAKE_INSTALL_PREFIX}/bin") + +# avoid overwrite of user defined settings +# install `config.txt`if file not exists in `${CMAKE_INSTALL_PREFIX}/bin` +install(CODE " \ + if(NOT EXISTS ${CMAKE_INSTALL_PREFIX}/bin/config.txt)\n \ + file(INSTALL ${CMAKE_CURRENT_SOURCE_DIR}/config.txt \ + DESTINATION ${CMAKE_INSTALL_PREFIX}/bin)\n \ + endif()" +) diff --git a/cli-miner.cpp b/cli-miner.cpp index 365ea27..5ef172b 100644 --- a/cli-miner.cpp +++ b/cli-miner.cpp @@ -26,7 +26,10 @@ #include "jconf.h" #include "console.h" #include "donate-level.h" -#include "httpd.h" + +#ifndef CONF_NO_HTTPD +# include "httpd.h" +#endif #include #include @@ -109,6 +112,7 @@ int main(int argc, char *argv[]) return 0; } +#ifndef CONF_NO_HTTPD if(jconf::inst()->GetHttpdPort() != 0) { if (!httpd::inst()->start_daemon()) @@ -117,6 +121,7 @@ int main(int argc, char *argv[]) return 0; } } +#endif printer::inst()->print_str("-------------------------------------------------------------------\n"); printer::inst()->print_str("XMR-Stak-AMD mining software, AMD Version.\n"); diff --git a/httpd.cpp b/httpd.cpp index 7e82aa3..7c94f76 100644 --- a/httpd.cpp +++ b/httpd.cpp @@ -21,6 +21,8 @@ * */ +#ifndef CONF_NO_HTTPD + #include #include #include @@ -142,3 +144,5 @@ bool httpd::start_daemon() return true; } +#endif + From 28b29179f4bb0e4ed36d366dcc98df00e6112944 Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Sun, 30 Apr 2017 19:54:23 +0200 Subject: [PATCH 04/16] always build feature complete - throw an error if a compile dependancy is not fullfilled - add options to disable hard compile dependancies --- CMakeLists.txt | 25 +++++++++++++++++++------ 1 file changed, 19 insertions(+), 6 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index acf4093..7a25a6f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -56,10 +56,15 @@ set(LIBS ${LIBS} ${CMAKE_THREAD_LIBS_INIT}) # Find microhttpd ################################################################################ +option(MICROHTTPD_REQUIRED "Enable or disable the requirement of microhttp (http deamon)" ON) find_library(MHTD NAMES microhttpd) if("${MHTD}" STREQUAL "MHTD-NOTFOUND") - message(STATUS "microhttpd NOT found: disable http server") - add_definitions("-DCONF_NO_HTTPD") + if(MICROHTTPD_REQUIRED) + message(FATAL_ERROR "microhttpd NOT found: use `-DMICROHTTPD_REQUIRED=OFF` to build without http deamon support") + else() + message(STATUS "microhttpd NOT found: disable http server") + add_definitions("-DCONF_NO_HTTPD") + endif() else() set(LIBS ${LIBS} ${MHTD}) endif() @@ -68,11 +73,19 @@ endif() # Find OpenSSL ############################################################################### +option(OpenSSL_REQUIRED "Enable or disable the requirement of OpenSSL" ON) find_package(OpenSSL) -include_directories(${OPENSSL_INCLUDE_DIR}) -set(LIBS ${LIBS} ${OPENSSL_LIBRARIES}) -if(NOT OPENSSL_FOUND) - add_definitions("-DCONF_NO_TLS") +if(OPENSSL_FOUND) + include_directories(${OPENSSL_INCLUDE_DIR}) + set(LIBS ${LIBS} ${OPENSSL_LIBRARIES}) +else() + if(OpenSSL_REQUIRED) + message(FATAL_ERROR "OpenSSL NOT found: use `-DOpenSSL_REQUIRED=OFF` to build without SSL support") + else() + if(NOT OPENSSL_FOUND) + add_definitions("-DCONF_NO_TLS") + endif() + endif() endif() ################################################################################ From bffbc9d3740053c99ac42481466021934b5beae6 Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Thu, 20 Apr 2017 21:21:10 +0200 Subject: [PATCH 05/16] allow optional static linking - add cmake option `CMAKE_LINK_STATIC` - add missing `CMAKE_C_FLAGS` --- CMakeLists.txt | 15 +++++++++++++-- 1 file changed, 13 insertions(+), 2 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index acf4093..f00483b 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -36,6 +36,9 @@ if(NOT CMAKE_BUILD_TYPE) endif() set_property(CACHE CMAKE_BUILD_TYPE PROPERTY STRINGS "${BUILD_TYPE}") +# option to add static libgcc and libstdc++ +option(CMAKE_LINK_STATIC "link as much as possible libraries static" OFF) + ############################################################################### # Find OpenCL ############################################################################### @@ -83,6 +86,15 @@ include_directories(.) # activate sse2 and aes-ni set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -msse2 -maes") +set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -msse2 -maes") + +# activate static libgcc and libstdc++ linking +if(CMAKE_LINK_STATIC) + set(BUILD_SHARED_LIBRARIES OFF) + set(DL_LIB ${CMAKE_DL_LIBS}) + set(CMAKE_FIND_LIBRARY_SUFFIXES ".a") + set(LIBS "-static-libgcc -static-libstdc++ ${LIBS}") +endif() file(GLOB SRCFILES_CPP "*.cpp" "crypto/*.cpp") file(GLOB SRCFILES_C "crypto/*.c" "amd_gpu/*.c") @@ -98,7 +110,7 @@ add_executable(xmr-stak-amd ) set(EXECUTABLE_OUTPUT_PATH "bin") -target_link_libraries(xmr-stak-amd xmr-stak-c ${LIBS}) +target_link_libraries(xmr-stak-amd ${LIBS} xmr-stak-c) ################################################################################ # Install @@ -121,4 +133,3 @@ install(CODE " \ DESTINATION ${CMAKE_INSTALL_PREFIX}/bin)\n \ endif()" ) - From 85f2b67df0d6d23ccf17ea6eb7827ab6afb79be0 Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Mon, 1 May 2017 21:11:22 +0200 Subject: [PATCH 06/16] allow to fully disable OpenSSL and microhttpd - rename option `*_REQUIRED` to `*_ENABLE` - allow to disable the dependency OpenSSL and microhttpd --- CMakeLists.txt | 35 ++++++++++++++++------------------- 1 file changed, 16 insertions(+), 19 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 7a25a6f..46606b3 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -56,36 +56,33 @@ set(LIBS ${LIBS} ${CMAKE_THREAD_LIBS_INIT}) # Find microhttpd ################################################################################ -option(MICROHTTPD_REQUIRED "Enable or disable the requirement of microhttp (http deamon)" ON) -find_library(MHTD NAMES microhttpd) -if("${MHTD}" STREQUAL "MHTD-NOTFOUND") - if(MICROHTTPD_REQUIRED) - message(FATAL_ERROR "microhttpd NOT found: use `-DMICROHTTPD_REQUIRED=OFF` to build without http deamon support") +option(MICROHTTPD_ENABLE "Enable or disable the requirement of microhttp (http deamon)" ON) +if(MICROHTTPD_ENABLE) + find_library(MHTD NAMES microhttpd) + if("${MHTD}" STREQUAL "MHTD-NOTFOUND") + message(FATAL_ERROR "microhttpd NOT found: use `-DMICROHTTPD_ENABLE=OFF` to build without http deamon support") else() - message(STATUS "microhttpd NOT found: disable http server") - add_definitions("-DCONF_NO_HTTPD") + set(LIBS ${LIBS} ${MHTD}) endif() else() - set(LIBS ${LIBS} ${MHTD}) + add_definitions("-DCONF_NO_HTTPD") endif() ############################################################################### # Find OpenSSL ############################################################################### -option(OpenSSL_REQUIRED "Enable or disable the requirement of OpenSSL" ON) -find_package(OpenSSL) -if(OPENSSL_FOUND) - include_directories(${OPENSSL_INCLUDE_DIR}) - set(LIBS ${LIBS} ${OPENSSL_LIBRARIES}) -else() - if(OpenSSL_REQUIRED) - message(FATAL_ERROR "OpenSSL NOT found: use `-DOpenSSL_REQUIRED=OFF` to build without SSL support") +option(OpenSSL_ENABLE "Enable or disable the requirement of OpenSSL" ON) +if(OpenSSL_ENABLE) + find_package(OpenSSL) + if(OPENSSL_FOUND) + include_directories(${OPENSSL_INCLUDE_DIR}) + set(LIBS ${LIBS} ${OPENSSL_LIBRARIES}) else() - if(NOT OPENSSL_FOUND) - add_definitions("-DCONF_NO_TLS") - endif() + message(FATAL_ERROR "OpenSSL NOT found: use `-DOpenSSL_ENABLE=OFF` to build without SSL support") endif() +else() + add_definitions("-DCONF_NO_TLS") endif() ################################################################################ From 486057a2f25a1dc7759b8145b4087ba1f1fb780f Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Mon, 1 May 2017 21:37:39 +0200 Subject: [PATCH 07/16] add file version.h - add file with version number - add version number to the terminal output --- cli-miner.cpp | 3 ++- jpsock.cpp | 3 ++- version.h | 4 ++++ 3 files changed, 8 insertions(+), 2 deletions(-) create mode 100644 version.h diff --git a/cli-miner.cpp b/cli-miner.cpp index 5ef172b..a4c509b 100644 --- a/cli-miner.cpp +++ b/cli-miner.cpp @@ -26,6 +26,7 @@ #include "jconf.h" #include "console.h" #include "donate-level.h" +#include "version.h" #ifndef CONF_NO_HTTPD # include "httpd.h" @@ -124,7 +125,7 @@ int main(int argc, char *argv[]) #endif printer::inst()->print_str("-------------------------------------------------------------------\n"); - printer::inst()->print_str("XMR-Stak-AMD mining software, AMD Version.\n"); + printer::inst()->print_str( XMR_STAK_NAME" " XMR_STAK_VERSION " mining software, AMD Version.\n"); printer::inst()->print_str("AMD mining code was written by wolf9466.\n"); printer::inst()->print_str("Brought to you by fireice_uk under GPLv3.\n\n"); char buffer[64]; diff --git a/jpsock.cpp b/jpsock.cpp index 70241f9..26af165 100644 --- a/jpsock.cpp +++ b/jpsock.cpp @@ -32,8 +32,9 @@ #include "jext.h" #include "socks.h" #include "socket.h" +#include "version.h" -#define AGENTID_STR "xmr-stak-amd/1.3.1" +#define AGENTID_STR XMR_STAK_NAME "/" XMR_STAK_VERSION using namespace rapidjson; diff --git a/version.h b/version.h new file mode 100644 index 0000000..b1d415d --- /dev/null +++ b/version.h @@ -0,0 +1,4 @@ +#pragma once + +#define XMR_STAK_NAME "xmr-stak-amd" +#define XMR_STAK_VERSION "1.3.1-dev" From 7d176e7d2fe288d329fd8cceb7998d81ff50f906 Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Mon, 15 May 2017 21:00:59 +0200 Subject: [PATCH 08/16] add missing include add include `math.h` --- amd_gpu/gpu.c | 1 + 1 file changed, 1 insertion(+) diff --git a/amd_gpu/gpu.c b/amd_gpu/gpu.c index 547b1cf..a888aef 100644 --- a/amd_gpu/gpu.c +++ b/amd_gpu/gpu.c @@ -15,6 +15,7 @@ #include #include +#include #ifdef _WIN32 #include From 6a3624927501dc5ffc4dbc2508872b9a601f4541 Mon Sep 17 00:00:00 2001 From: Cian Montgomery Date: Sun, 21 May 2017 09:54:24 -0700 Subject: [PATCH 09/16] Fix cmake build on linux, OpenCL link dependencies missing from C library. Signed-off-by: Cian Montgomery --- CMakeLists.txt | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 575d104..de75f39 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -109,18 +109,18 @@ endif() file(GLOB SRCFILES_CPP "*.cpp" "crypto/*.cpp") file(GLOB SRCFILES_C "crypto/*.c" "amd_gpu/*.c") -add_library(xmr-stak-c +add_library(xmr-stak-amd-c STATIC ${SRCFILES_C} ) -set_property(TARGET xmr-stak-c PROPERTY C_STANDARD 99) +set_property(TARGET xmr-stak-amd-c PROPERTY C_STANDARD 99) +target_link_libraries(xmr-stak-amd-c PUBLIC ${OpenCL_LIBRARY}) add_executable(xmr-stak-amd ${SRCFILES_CPP} ) - set(EXECUTABLE_OUTPUT_PATH "bin") -target_link_libraries(xmr-stak-amd ${LIBS} xmr-stak-c) +target_link_libraries(xmr-stak-amd ${LIBS} xmr-stak-amd-c) ################################################################################ # Install From af01b2f16d55b488a9fae268f21400fe580141d1 Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Sun, 21 May 2017 21:45:44 +0200 Subject: [PATCH 10/16] optimize scratchpad memory layout store one scatchpad hash as block to optimize cash hits --- opencl/cryptonight.cl | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) diff --git a/opencl/cryptonight.cl b/opencl/cryptonight.cl index 764c3af..42a3ee2 100644 --- a/opencl/cryptonight.cl +++ b/opencl/cryptonight.cl @@ -346,7 +346,7 @@ void AESExpandKey256(uint *keybuf) } } -#define IDX(x) ((x) * (get_global_size(0))) +#define IDX(x) (x) __attribute__((reqd_work_group_size(WORKSIZE, 8, 1))) __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states) @@ -357,7 +357,7 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul uint4 text; states += (25 * (get_global_id(0) - get_global_offset(0))); - Scratchpad += ((get_global_id(0) - get_global_offset(0))); + Scratchpad += ((get_global_id(0) - get_global_offset(0))) * (0x80000 >> 2); for(int i = get_local_id(0); i < 256; i += WORKSIZE) { @@ -367,6 +367,7 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul AES2[i] = rotate(tmp, 16U); AES3[i] = rotate(tmp, 24U); } + barrier(CLK_LOCAL_MEM_FENCE); ((ulong8 *)State)[0] = vload8(0, input); State[8] = input[8]; @@ -418,7 +419,7 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states) ulong a[2], b[2]; __local uint AES0[256], AES1[256], AES2[256], AES3[256]; - Scratchpad += ((get_global_id(0) - get_global_offset(0))); + Scratchpad += ((get_global_id(0) - get_global_offset(0))) * (0x80000 >> 2); states += (25 * (get_global_id(0) - get_global_offset(0))); for(int i = get_local_id(0); i < 256; i += WORKSIZE) @@ -429,6 +430,7 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states) AES2[i] = rotate(tmp, 16U); AES3[i] = rotate(tmp, 24U); } + barrier(CLK_LOCAL_MEM_FENCE); a[0] = states[0] ^ states[4]; b[0] = states[2] ^ states[6]; @@ -474,7 +476,7 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u ulong State[25]; uint4 text; - Scratchpad += ((get_global_id(0) - get_global_offset(0))); + Scratchpad += ((get_global_id(0) - get_global_offset(0))) * (0x80000 >> 2); states += (25 * (get_global_id(0) - get_global_offset(0))); for(int i = get_local_id(0); i < 256; i += WORKSIZE) @@ -485,6 +487,7 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u AES2[i] = rotate(tmp, 16U); AES3[i] = rotate(tmp, 24U); } + barrier(CLK_LOCAL_MEM_FENCE); #if defined(__Tahiti__) || defined(__Pitcairn__) From 2b1d841753e92c390ae3879f596e047c8e684735 Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Fri, 26 May 2017 20:33:39 +0200 Subject: [PATCH 11/16] set version to 1.4.0 --- version.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/version.h b/version.h index b1d415d..1ee1381 100644 --- a/version.h +++ b/version.h @@ -1,4 +1,4 @@ #pragma once #define XMR_STAK_NAME "xmr-stak-amd" -#define XMR_STAK_VERSION "1.3.1-dev" +#define XMR_STAK_VERSION "1.4.0" From a6e6dd2940c8f70ba6ec700b6306c6c8d7eb4136 Mon Sep 17 00:00:00 2001 From: fireice-uk Date: Fri, 26 May 2017 18:59:49 +0100 Subject: [PATCH 12/16] Add the get_key() guard from the other branch --- cli-miner.cpp | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/cli-miner.cpp b/cli-miner.cpp index a4c509b..051a00f 100644 --- a/cli-miner.cpp +++ b/cli-miner.cpp @@ -142,6 +142,9 @@ int main(int argc, char *argv[]) executor::inst()->ex_start(); + using namespace std::chrono; + uint64_t lastTime = time_point_cast(high_resolution_clock::now()).time_since_epoch().count(); + int key; while(true) { @@ -161,6 +164,13 @@ int main(int argc, char *argv[]) default: break; } + + uint64_t currentTime = time_point_cast(high_resolution_clock::now()).time_since_epoch().count(); + + /* Hard guard to make sure we never get called more than twice per second */ + if( currentTime - lastTime < 500) + std::this_thread::sleep_for(std::chrono::milliseconds(500 - (currentTime - lastTime))); + lastTime = currentTime; } return 0; From 654d82d1253377d88052176a371fbc4f0571b0c3 Mon Sep 17 00:00:00 2001 From: fireice-uk Date: Fri, 26 May 2017 19:36:38 +0100 Subject: [PATCH 13/16] daemon mode feature --- cli-miner.cpp | 2 +- config.txt | 8 ++++++++ executor.cpp | 1 - executor.h | 6 +++--- jconf.cpp | 8 +++++++- jconf.h | 2 ++ 6 files changed, 21 insertions(+), 6 deletions(-) diff --git a/cli-miner.cpp b/cli-miner.cpp index 051a00f..252d888 100644 --- a/cli-miner.cpp +++ b/cli-miner.cpp @@ -140,7 +140,7 @@ int main(int argc, char *argv[]) if(strlen(jconf::inst()->GetOutputFile()) != 0) printer::inst()->open_logfile(jconf::inst()->GetOutputFile()); - executor::inst()->ex_start(); + executor::inst()->ex_start(jconf::inst()->DaemonMode()); using namespace std::chrono; uint64_t lastTime = time_point_cast(high_resolution_clock::now()).time_since_epoch().count(); diff --git a/config.txt b/config.txt index 0d6e254..89ccc7b 100644 --- a/config.txt +++ b/config.txt @@ -87,6 +87,14 @@ */ "h_print_time" : 60, +/* + * Daemon mode + * + * If you are running the process in the background and you don't need the keyboard reports, set this to true. + * This should solve the hashrate problems on some emulated terminals. + */ +"daemon_mode" : false, + /* * Output file * diff --git a/executor.cpp b/executor.cpp index 6d35f1b..149725e 100644 --- a/executor.cpp +++ b/executor.cpp @@ -49,7 +49,6 @@ executor* executor::oInst = NULL; executor::executor() { - my_thd = nullptr; cpu_ctx = (cryptonight_ctx*)_mm_malloc(sizeof(cryptonight_ctx), 4096); } diff --git a/executor.h b/executor.h index a6c3444..e93a5b4 100644 --- a/executor.h +++ b/executor.h @@ -21,8 +21,7 @@ public: return oInst; }; - void ex_start() { my_thd = new std::thread(&executor::ex_main, this); } - void ex_main(); + void ex_start(bool daemon) { daemon ? ex_main() : std::thread(&executor::ex_main, this).detach(); } void get_http_report(ex_event_name ev_id, std::string& data); @@ -57,7 +56,6 @@ private: telemetry* telem; std::vector* pvThreads; - std::thread* my_thd; size_t current_pool_id; @@ -71,6 +69,8 @@ private: executor(); static executor* oInst; + void ex_main(); + void ex_clock_thd(); void pool_connect(jpsock* pool); diff --git a/jconf.cpp b/jconf.cpp index 9777abe..cae79eb 100644 --- a/jconf.cpp +++ b/jconf.cpp @@ -48,7 +48,7 @@ using namespace rapidjson; enum configEnum { iGpuThreadNum, aGpuThreadsConf, iPlatformIdx, bTlsMode, bTlsSecureAlgo, sTlsFingerprint, sPoolAddr, sWalletAddr, sPoolPwd, iCallTimeout, iNetRetry, iGiveUpLimit, iVerboseLevel, iAutohashTime, - sOutputFile, iHttpdPort, bPreferIpv4 }; + bDaemonMode, sOutputFile, iHttpdPort, bPreferIpv4 }; struct configVal { configEnum iName; @@ -72,6 +72,7 @@ configVal oConfigValues[] = { { iGiveUpLimit, "giveup_limit", kNumberType }, { iVerboseLevel, "verbose_level", kNumberType }, { iAutohashTime, "h_print_time", kNumberType }, + { bDaemonMode, "daemon_mode", kTrueType }, { sOutputFile, "output_file", kStringType }, { iHttpdPort, "httpd_port", kNumberType }, { bPreferIpv4, "prefer_ipv4", kTrueType } @@ -220,6 +221,11 @@ uint16_t jconf::GetHttpdPort() return prv->configValues[iHttpdPort]->GetUint(); } +bool jconf::DaemonMode() +{ + return prv->configValues[bDaemonMode]->GetBool(); +} + const char* jconf::GetOutputFile() { return prv->configValues[sOutputFile]->GetString(); diff --git a/jconf.h b/jconf.h index 51af4d2..36ef0de 100644 --- a/jconf.h +++ b/jconf.h @@ -44,6 +44,8 @@ public: uint16_t GetHttpdPort(); + bool DaemonMode(); + bool PreferIpv4(); inline bool HaveHardwareAes() { return bHaveAes; } From 3e2927a2472200e816fb19bfcea9a2ea3eab647e Mon Sep 17 00:00:00 2001 From: fireice-uk Date: Sun, 28 May 2017 09:37:24 +0100 Subject: [PATCH 14/16] version --- version.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/version.h b/version.h index 1ee1381..57d5ff9 100644 --- a/version.h +++ b/version.h @@ -1,4 +1,4 @@ #pragma once #define XMR_STAK_NAME "xmr-stak-amd" -#define XMR_STAK_VERSION "1.4.0" +#define XMR_STAK_VERSION "1.1.0-1.4.0" From 93a92978d2d49feebb0f0a9d8469fb1ac2c98414 Mon Sep 17 00:00:00 2001 From: fireice-uk Date: Sun, 28 May 2017 10:20:11 +0100 Subject: [PATCH 15/16] Update checksums --- README.md | 28 ++++++++++++++-------------- 1 file changed, 14 insertions(+), 14 deletions(-) diff --git a/README.md b/README.md index 47dac88..475180f 100644 --- a/README.md +++ b/README.md @@ -21,41 +21,41 @@ Hash: SHA256 Windows binary release checksums sha1sum -d45ca7cbab2ca9fe994a6c8749be08cc9a13ba69 xmr-stak-amd.exe -4892d4fcb6ef6f6132a90bbb628549eb4b5dc5fd xmr-stak-amd-notls.exe +774f20271610741ec20c5f1c7e9d20467e3f3834 xmr-stak-amd.exe +cc7c22b0800a4c615d05dd53f8c6faa65d20ba24 xmr-stak-amd-notls.exe d34a0ba0dd7b3b1f900a7e02772e197e974b4a73 libeay32.dll 2ee9966a0fc163da58408d91be36b84fa287c10b ssleay32.dll 5acb656005a86cad90c6327985c3795d96a96c91 opencl/blake256.cl -9e4e276afd9000945c25f6e5a5259977a29239f4 opencl/cryptonight.cl +603322cbe721964d76de68660f63f97f1bcd4057 opencl/cryptonight.cl c9fb5e4bfb137ff60063978eecd10bffb7c4deb6 opencl/groestl256.cl 361dfce776ee4a89b100d139a879af5d0f0d65e4 opencl/jh.cl 429f559190d1163335847cc08df955234051504b opencl/wolf-aes.cl e2862a6d7094aeab21844d3155803e5da99b4a46 opencl/wolf-skein.cl sha3sum -b5749e46c07793c39a651f41ca8a617f4c3ab2072c50b208734afc6f xmr-stak-amd.exe -fb29fc2059af15d0bbcd9cb73382d226c7023c22219a6e7edaa07ffc xmr-stak-amd-notls.exe +487b2f05cfd71a9cb4acd0eba2734162d605ada539900f83f43d7a6b xmr-stak-amd.exe +d79dec89112811ea61a4486a91e1c22e48f3293b1d6e28e951d346af xmr-stak-amd-notls.exe 133c065d9ef2c93396382e2ba5d8c3ca8c6a57c6beb0159cb9a4b6c5 ssleay32.dll 05003137a87313c81d6c348c9b96411c95d48dc22c35f36c39129747 libeay32.dll a52b05548fd094e7bbb2367d7922bf19af3ed23ad5df53004fae0825 opencl/blake256.cl -a49c9da554d7b01d091eea56e8b97b943ca33cd2a64a1f3f3169e202 opencl/cryptonight.cl +8b3d0f4a39fca8a6d21afe9121796343379605017a167fa148c374d3 opencl/cryptonight.cl 13315b0a475212c92e10b7627b03a0813132437d4496b6596821b909 opencl/groestl256.cl 89548b917dbe216f5809624ebe350859c1d800048909322049f93d23 opencl/jh.cl 806eb1d4e3d7b6630a422bb42ee00faa76d31143b7c1cbde65e46938 opencl/wolf-aes.cl 052176a740a5a0bc088feea0aa7a72f0e9d96d6b6ffd00844676dd17 opencl/wolf-skein.cl date -Wed 15 Mar 16:43:43 GMT 2017 +Sun 28 May 10:18:31 BST 2017 -----BEGIN PGP SIGNATURE----- Version: GnuPG v2 -iQEcBAEBCAAGBQJYyW9uAAoJEPsk95p+1Bw0olMIAKCd6SbJqj6mrX4u8GtYAF5X -q0OJc/B27P9BRmP5bjPCJg48iopwe2MYHJN+We6ygw6+y/gn9VNdPnznoxgCaRRp -MEZzg4UYcz3zBlDe5Vfs56PbChky5x9sg02rPzVAX84SakNhRkOcS3638GDL5c20 -4gldasMOHyqnM+mjlM8RpOev3SR18hXJAjGam3XxGTFqaI70WdjlwSpCwoQ9Cm6I -BIBsrTn7CK9gYn6RGj71ZDc7azW8GMPe+Aq7X3NP9VmkjmtrkjOrlwECQC20KQ6F -kSRwQQKvgsqFcxoFbmvx59X61qg4amfgEe4VLb8+lh6ZBs5TVRuzJmLZnn2v3dM= -=xYQ1 +iQEcBAEBCAAGBQJZKpX9AAoJEPsk95p+1Bw0BAYIALFYF5Ud+kKxb9uNyQaD3h5n +yv1MrdI81lpTRWpjK2naAN+pymAx7MwMi4KU0hTy5FMOuX+Ssb8kJsCAXKUYogff +W4+M4Bk6JaCcrJVhhgK2ucuDm8H1uQ+Ps/nPa+a7foT6d+7kGtj7bXPeOWyr+oUB +IbP/LnG40jOJnzG1u6/1e3YLXewWOEdLOYi8fN7VAl5b5uwmWw4DcWkQT1aiaFCV +awJCjDUNWtZ3+wgzE9ZU2cPmNXDKTuVHnR9F6+S8CDJs+yPe5phQuTrIObB5dhpD +trZnODWL/fZehf1koNCxKxpBDhGaDX7nx7j0NjN/btYhNwm9dckd9Xz3B3/Cod8= +=0alz -----END PGP SIGNATURE----- ``` From 70d615c41833f2144f01543e5805c754591608fe Mon Sep 17 00:00:00 2001 From: fireice-uk Date: Sun, 28 May 2017 10:27:32 +0100 Subject: [PATCH 16/16] 32+ cpu fix, already rolled in hashed binary --- minethd.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/minethd.cpp b/minethd.cpp index 02fe275..728abbb 100644 --- a/minethd.cpp +++ b/minethd.cpp @@ -32,7 +32,7 @@ void thd_setaffinity(std::thread::native_handle_type h, uint64_t cpu_id) { - SetThreadAffinityMask(h, 1 << cpu_id); + SetThreadAffinityMask(h, 1ULL << cpu_id); } #else