fireice-uk
GitHub
commit
48e358a46d
145
CMakeLists.txt
145
CMakeLists.txt
|
|
@ -1,38 +1,145 @@
|
|||
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()
|
||||
|
||||
find_library(MHTD NAMES microhttpd)
|
||||
if("${MHTD}" STREQUAL "MHTD-NOTFOUND")
|
||||
message(FATAL_ERROR "libmicrohttpd is required")
|
||||
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}")
|
||||
|
||||
# option to add static libgcc and libstdc++
|
||||
option(CMAKE_LINK_STATIC "link as much as possible libraries static" OFF)
|
||||
|
||||
###############################################################################
|
||||
# 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
|
||||
################################################################################
|
||||
|
||||
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()
|
||||
set(LIBS ${LIBS} ${MHTD})
|
||||
endif()
|
||||
else()
|
||||
add_definitions("-DCONF_NO_HTTPD")
|
||||
endif()
|
||||
|
||||
find_package(OpenSSL REQUIRED)
|
||||
include_directories(${OPENSSL_INCLUDE_DIR})
|
||||
###############################################################################
|
||||
# Find OpenSSL
|
||||
###############################################################################
|
||||
|
||||
#set(CMAKE_VERBOSE_MAKEFILE ON)
|
||||
set(CMAKE_CONFIGURATION_TYPES "RELEASE;STATIC")
|
||||
if("${CMAKE_BUILD_TYPE}" STREQUAL "")
|
||||
set(CMAKE_BUILD_TYPE RELEASE)
|
||||
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()
|
||||
message(FATAL_ERROR "OpenSSL NOT found: use `-DOpenSSL_ENABLE=OFF` to build without SSL support")
|
||||
endif()
|
||||
else()
|
||||
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")
|
||||
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")
|
||||
|
||||
add_library(xmr-stak-amd-c
|
||||
STATIC
|
||||
${SRCFILES_C}
|
||||
)
|
||||
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-amd-c)
|
||||
|
||||
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()"
|
||||
)
|
||||
|
|
|
|||
28
README.md
28
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-----
|
||||
```
|
||||
|
||||
|
|
|
|||
|
|
@ -15,6 +15,7 @@
|
|||
|
||||
#include <stdio.h>
|
||||
#include <string.h>
|
||||
#include <math.h>
|
||||
|
||||
#ifdef _WIN32
|
||||
#include <windows.h>
|
||||
|
|
@ -163,10 +164,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";
|
||||
}
|
||||
|
|
@ -675,7 +678,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);
|
||||
|
|
|
|||
|
|
@ -26,7 +26,11 @@
|
|||
#include "jconf.h"
|
||||
#include "console.h"
|
||||
#include "donate-level.h"
|
||||
#include "httpd.h"
|
||||
#include "version.h"
|
||||
|
||||
#ifndef CONF_NO_HTTPD
|
||||
# include "httpd.h"
|
||||
#endif
|
||||
|
||||
#include <stdlib.h>
|
||||
#include <stdio.h>
|
||||
|
|
@ -109,6 +113,7 @@ int main(int argc, char *argv[])
|
|||
return 0;
|
||||
}
|
||||
|
||||
#ifndef CONF_NO_HTTPD
|
||||
if(jconf::inst()->GetHttpdPort() != 0)
|
||||
{
|
||||
if (!httpd::inst()->start_daemon())
|
||||
|
|
@ -117,9 +122,10 @@ 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");
|
||||
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];
|
||||
|
|
@ -134,7 +140,10 @@ 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<milliseconds>(high_resolution_clock::now()).time_since_epoch().count();
|
||||
|
||||
int key;
|
||||
while(true)
|
||||
|
|
@ -155,6 +164,13 @@ int main(int argc, char *argv[])
|
|||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
uint64_t currentTime = time_point_cast<milliseconds>(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;
|
||||
|
|
|
|||
|
|
@ -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
|
||||
*
|
||||
|
|
|
|||
|
|
@ -49,7 +49,6 @@ executor* executor::oInst = NULL;
|
|||
|
||||
executor::executor()
|
||||
{
|
||||
my_thd = nullptr;
|
||||
cpu_ctx = (cryptonight_ctx*)_mm_malloc(sizeof(cryptonight_ctx), 4096);
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -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<minethd*>* 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);
|
||||
|
||||
|
|
|
|||
|
|
@ -21,6 +21,8 @@
|
|||
*
|
||||
*/
|
||||
|
||||
#ifndef CONF_NO_HTTPD
|
||||
|
||||
#include <stdlib.h>
|
||||
#include <stdio.h>
|
||||
#include <string.h>
|
||||
|
|
@ -142,3 +144,5 @@ bool httpd::start_daemon()
|
|||
return true;
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
|
|
|
|||
|
|
@ -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();
|
||||
|
|
|
|||
2
jconf.h
2
jconf.h
|
|
@ -44,6 +44,8 @@ public:
|
|||
|
||||
uint16_t GetHttpdPort();
|
||||
|
||||
bool DaemonMode();
|
||||
|
||||
bool PreferIpv4();
|
||||
|
||||
inline bool HaveHardwareAes() { return bHaveAes; }
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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__)
|
||||
|
||||
|
|
|
|||
Loading…
Reference in New Issue