diff options
| -rw-r--r-- | README.md | 2 | ||||
| -rw-r--r-- | doc/FAQ.md | 19 | ||||
| -rw-r--r-- | doc/Linux_deployment.md | 1 | ||||
| -rw-r--r-- | doc/compile_Linux.md | 28 | ||||
| -rw-r--r-- | xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl | 104 | ||||
| -rw-r--r-- | xmrstak/backend/amd/minethd.cpp | 2 | ||||
| -rw-r--r-- | xmrstak/backend/cpu/crypto/cryptonight_aesni.h | 13 | ||||
| -rw-r--r-- | xmrstak/backend/nvidia/minethd.cpp | 2 | ||||
| -rw-r--r-- | xmrstak/misc/executor.cpp | 2 | ||||
| -rw-r--r-- | xmrstak/misc/uac.cpp | 2 | ||||
| -rw-r--r-- | xmrstak/net/msgstruct.hpp | 5 |
11 files changed, 100 insertions, 80 deletions
@@ -62,7 +62,7 @@ fireice-uk: psychocrypt: ``` -43NoJVEXo21hGZ6tDG6Z3g4qimiGdJPE6GRxAmiWwm26gwr62Lqo7zRiCJFSBmbkwTGNuuES9ES5TgaVHceuYc4Y75txCTU +45tcqnJMgd3VqeTznNotiNj4G9PQoK67TGRiHyj6EYSZ31NUbAfs9XdiU5squmZb717iHJLxZv3KfEw8jCYGL5wa19yrVCn ``` ## Release Checksums @@ -1,7 +1,7 @@ # FAQ ## Content Overview -* [SeLockMemoryPrivilege failed](#selockmemoryprivilege-failed) +* ["Obtaining SeLockMemoryPrivilege failed."](#obtaining-selockmemoryprivilege-failed) * [VirtualAlloc failed](#virtualalloc-failed) * [Error msvcp140.dll and vcruntime140.dll not available](#error-msvcp140dll-and-vcruntime140dll-not-available) * [Error: MEMORY ALLOC FAILED: mmap failed](#error-memory-alloc-failed-mmap-failed) @@ -9,23 +9,24 @@ * [Virus Protection Alert](#virus-protection-alert) * [Change Currency to Mine](#change-currency-to-mine) -## SeLockMemoryPrivilege failed +## "Obtaining SeLockMemoryPrivilege failed." -Please see [config.txt](config.txt) under section **LARGE PAGE SUPPORT** +For professional versions of Windows see [this article](https://msdn.microsoft.com/en-gb/library/ms190730.aspx). +Make sure to reboot afterwards! -For Windows 7 pro, or Windows 8 and above see [this article](https://msdn.microsoft.com/en-gb/library/ms190730.aspx) (make sure to reboot afterwards!). +For Windows 7/10 Home: -For Windows 7 Home : +1) Download and install [Windows Server 2003 Resource Kit Tools](https://www.microsoft.com/en-us/download/details.aspx?id=17657). Ignore any incompatibility warning during installation. -1) Download and install [Windows Server 2003 Resource Kit Tools](https://www.microsoft.com/en-us/download/details.aspx?id=17657). Ignore incompatiablity warning during installation. +2) Open cmd or PowerShell as an administrator. -2) In cmd or power shell: `ntrights -u %USERNAME% +r SeLockMemoryPrivilege` (where %USERNAME% is the user that will be running the program. This command needs to be run as admin) +3) Use `ntrights -u %USERNAME% +r SeLockMemoryPrivilege` where %USERNAME% is the user that will be running the program. -3) Reboot. +4) Reboot. Reference: http://rybkaforum.net/cgi-bin/rybkaforum/topic_show.pl?pid=259791#pid259791 -*Warning: do not download ntrights.exe from any other site other then the offical Microsoft download page.* +*Warning: Do not download ntrights.exe from any other site other than the offical Microsoft download page.* ## VirtualAlloc failed diff --git a/doc/Linux_deployment.md b/doc/Linux_deployment.md index 323a97f..3219e8a 100644 --- a/doc/Linux_deployment.md +++ b/doc/Linux_deployment.md @@ -15,6 +15,7 @@ For automatic deployments, please use the steps above to obtain config.txt and u ``` #!/bin/bash +sudo apt install curl curl -O `curl -s https://api.github.com/repos/fireice-uk/xmr-stak/releases/latest | grep -o 'browser_download_url.*xmr-stak-portbin-linux.tar.gz' | sed 's/.*"//'` curl -O http://path.to/your/config.txt tar xzf xmr-stak-portbin-linux.tar.gz diff --git a/doc/compile_Linux.md b/doc/compile_Linux.md index b7104ac..5a1e762 100644 --- a/doc/compile_Linux.md +++ b/doc/compile_Linux.md @@ -64,6 +64,34 @@ cd xmr-stak/build cmake .. make install + + # TinyCore Linux 8.x + # TinyCore is 32-bit only, but there is an x86-64 port, known as "Pure 64," + # hosted on the TinyCore home page, and it works well. + # Beware that huge page support is not enabled in the kernel distributed + # with Pure 64. Consider http://wiki.tinycorelinux.net/wiki:custom_kernel + # Note that as of yet there are no distro packages for microhttpd or hwloc. + # hwloc is easy enough to install manually though, shown below. + # Also note that only CPU mining has been tested on this platform, thus the + # disabling of CUDA and OpenCL shown below. + tce-load -iw openssl-dev.tcz cmake.tcz make.tcz gcc.tcz git.tcz \ + glibc_base-dev.tcz linux-4.8.1_api_headers.tcz \ + glibc_add_lib.tcz + wget https://www.open-mpi.org/software/hwloc/v1.11/downloads/hwloc-1.11.8.tar.gz + tar xzvf hwloc-1.11.8.tar.gz + cd hwloc-1.11.8 + ./configure --prefix=/usr/local + make + sudo make install + cd .. + git clone http://github.com/fireice-uk/xmr-stak + cd xmr-stak + mkdir build + cd build + CC=gcc cmake .. -DCUDA_ENABLE=OFF \ + -DOpenCL_ENABLE=OFF \ + -DMICROHTTPD_ENABLE=OFF + make install ``` - g++ version 5.1 or higher is required for full C++11 support. diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl index 255fcbb..ec05712 100644 --- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl +++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl @@ -653,21 +653,11 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u for(int i = 0; i < 25; ++i) states[i] = State[i]; - switch(State[0] & 3) - { - case 0: - Branch0[atomic_inc(Branch0 + Threads)] = get_global_id(0) - get_global_offset(0); - break; - case 1: - Branch1[atomic_inc(Branch1 + Threads)] = get_global_id(0) - get_global_offset(0); - break; - case 2: - Branch2[atomic_inc(Branch2 + Threads)] = get_global_id(0) - get_global_offset(0); - break; - case 3: - Branch3[atomic_inc(Branch3 + Threads)] = get_global_id(0) - get_global_offset(0); - break; - } + ulong StateSwitch = State[0] & 3; + __global uint *destinationBranch1 = StateSwitch == 0 ? Branch0 : Branch1; + __global uint *destinationBranch2 = StateSwitch == 2 ? Branch2 : Branch3; + __global uint *destinationBranch = StateSwitch < 2 ? destinationBranch1 : destinationBranch2; + destinationBranch[atomic_inc(destinationBranch + Threads)] = gIdx; } } mem_fence(CLK_GLOBAL_MEM_FENCE); @@ -704,8 +694,7 @@ __kernel void Skein(__global ulong *states, __global uint *BranchBuf, __global u for(uint i = 0; i < 4; ++i) { - if(i < 3) t[0] += 0x40UL; - else t[0] += 0x08UL; + t[0] += i < 3 ? 0x40UL : 0x08UL; t[2] = t[0] ^ t[1]; @@ -715,8 +704,7 @@ __kernel void Skein(__global ulong *states, __global uint *BranchBuf, __global u h = m ^ p; - if(i < 2) t[1] = 0x3000000000000000UL; - else t[1] = 0xB000000000000000UL; + t[1] = i < 2 ? 0x3000000000000000UL : 0xB000000000000000UL; } t[0] = 0x08UL; @@ -744,6 +732,27 @@ __kernel void Skein(__global ulong *states, __global uint *BranchBuf, __global u #define SWAP8(x) as_ulong(as_uchar8(x).s76543210) +#define JHXOR \ + h0h ^= input[0]; \ + h0l ^= input[1]; \ + h1h ^= input[2]; \ + h1l ^= input[3]; \ + h2h ^= input[4]; \ + h2l ^= input[5]; \ + h3h ^= input[6]; \ + h3l ^= input[7]; \ +\ + E8; \ +\ + h4h ^= input[0]; \ + h4l ^= input[1]; \ + h5h ^= input[2]; \ + h5l ^= input[3]; \ + h6h ^= input[4]; \ + h6l ^= input[5]; \ + h7h ^= input[6]; \ + h7l ^= input[7] + __kernel void JH(__global ulong *states, __global uint *BranchBuf, __global uint *output, ulong Target, ulong Threads) { const uint idx = get_global_id(0) - get_global_offset(0); @@ -757,46 +766,27 @@ __kernel void JH(__global ulong *states, __global uint *BranchBuf, __global uint sph_u64 h4h = 0x754D2E7F8996A371UL, h4l = 0x62E27DF70849141DUL, h5h = 0x948F2476F7957627UL, h5l = 0x6C29804757B6D587UL, h6h = 0x6C0D8EAC2D275E5CUL, h6l = 0x0F7A0557C6508451UL, h7h = 0xEA12247067D3E47BUL, h7l = 0x69D71CD313ABE389UL; sph_u64 tmp; - for(int i = 0; i < 5; ++i) + for(int i = 0; i < 3; ++i) { ulong input[8]; - if(i < 3) - { - for(int x = 0; x < 8; ++x) input[x] = (states[(i << 3) + x]); - } - else if(i == 3) - { - input[0] = (states[24]); - input[1] = 0x80UL; - for(int x = 2; x < 8; ++x) input[x] = 0x00UL; - } - else - { - input[7] = 0x4006000000000000UL; - - for(int x = 0; x < 7; ++x) input[x] = 0x00UL; - } - - h0h ^= input[0]; - h0l ^= input[1]; - h1h ^= input[2]; - h1l ^= input[3]; - h2h ^= input[4]; - h2l ^= input[5]; - h3h ^= input[6]; - h3l ^= input[7]; - - E8; - - h4h ^= input[0]; - h4l ^= input[1]; - h5h ^= input[2]; - h5l ^= input[3]; - h6h ^= input[4]; - h6l ^= input[5]; - h7h ^= input[6]; - h7l ^= input[7]; + const int shifted = i << 3; + for(int x = 0; x < 8; ++x) input[x] = (states[shifted + x]); + JHXOR; + } + { + ulong input[8]; + input[0] = (states[24]); + input[1] = 0x80UL; + #pragma unroll 6 + for(int x = 2; x < 8; ++x) input[x] = 0x00UL; + JHXOR; + } + { + ulong input[8]; + for(int x = 0; x < 7; ++x) input[x] = 0x00UL; + input[7] = 0x4006000000000000UL; + JHXOR; } //output[0] = h6h; @@ -832,6 +822,7 @@ __kernel void Blake(__global ulong *states, __global uint *BranchBuf, __global u ((uint8 *)h)[0] = vload8(0U, c_IV256); + #pragma unroll 4 for(uint i = 0, bitlen = 0; i < 4; ++i) { if(i < 3) @@ -907,6 +898,7 @@ __kernel void Groestl(__global ulong *states, __global uint *BranchBuf, __global State[7] = 0x0001000000000000UL; + #pragma unroll 4 for(uint i = 0; i < 4; ++i) { ulong H[8], M[8]; diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp index 0ee3f8e..422c28c 100644 --- a/xmrstak/backend/amd/minethd.cpp +++ b/xmrstak/backend/amd/minethd.cpp @@ -245,7 +245,7 @@ void minethd::work_main() if ( (*((uint64_t*)(bResult + 24))) < oWork.iTarget) executor::inst()->push_event(ex_event(job_result(oWork.sJobID, results[i], bResult, iThreadNo), oWork.iPoolId)); else - executor::inst()->push_event(ex_event("AMD Invalid Result", oWork.iPoolId)); + executor::inst()->push_event(ex_event("AMD Invalid Result", pGpuCtx->deviceIdx, oWork.iPoolId)); } iCount += pGpuCtx->rawIntensity; diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h index 9b6e1dc..e4ccbc3 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h +++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h @@ -317,10 +317,9 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c _mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx)); idx0 = _mm_cvtsi128_si64(cx); - bx0 = cx; - if(PREFETCH) _mm_prefetch((const char*)&l0[idx0 & MASK], _MM_HINT_T0); + bx0 = cx; uint64_t hi, lo, cl, ch; cl = ((uint64_t*)&l0[idx0 & MASK])[0]; @@ -329,15 +328,15 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c lo = _umul128(idx0, cl, &hi); al0 += hi; - ah0 += lo; ((uint64_t*)&l0[idx0 & MASK])[0] = al0; + al0 ^= cl; + if(PREFETCH) + _mm_prefetch((const char*)&l0[al0 & MASK], _MM_HINT_T0); + ah0 += lo; ((uint64_t*)&l0[idx0 & MASK])[1] = ah0; ah0 ^= ch; - al0 ^= cl; - idx0 = al0; - if(PREFETCH) - _mm_prefetch((const char*)&l0[idx0 & MASK], _MM_HINT_T0); + idx0 = al0; } // Optim - 90% time boundary diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp index cc6ea24..9fd08fb 100644 --- a/xmrstak/backend/nvidia/minethd.cpp +++ b/xmrstak/backend/nvidia/minethd.cpp @@ -287,7 +287,7 @@ void minethd::work_main() if ( (*((uint64_t*)(bResult + 24))) < oWork.iTarget) executor::inst()->push_event(ex_event(job_result(oWork.sJobID, foundNonce[i], bResult, iThreadNo), oWork.iPoolId)); else - executor::inst()->push_event(ex_event("NVIDIA Invalid Result", oWork.iPoolId)); + executor::inst()->push_event(ex_event("NVIDIA Invalid Result", ctx.device_id, oWork.iPoolId)); } iCount += h_per_round; diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp index a3088a5..f0b6e0e 100644 --- a/xmrstak/misc/executor.cpp +++ b/xmrstak/misc/executor.cpp @@ -594,7 +594,7 @@ void executor::ex_main() break; case EV_GPU_RES_ERROR: - log_result_error(std::string(ev.oGpuError.error_str)); + log_result_error(std::string(ev.oGpuError.error_str + std::string(" GPU ID ") + std::to_string(ev.oGpuError.idx))); break; case EV_PERF_TICK: diff --git a/xmrstak/misc/uac.cpp b/xmrstak/misc/uac.cpp index 5e8d08a..ad9d394 100644 --- a/xmrstak/misc/uac.cpp +++ b/xmrstak/misc/uac.cpp @@ -1,5 +1,3 @@ -#pragma once - #ifdef _WIN32 #include "xmrstak/misc/console.hpp" #include "xmrstak/params.hpp" diff --git a/xmrstak/net/msgstruct.hpp b/xmrstak/net/msgstruct.hpp index a5affc8..8c4bdbe 100644 --- a/xmrstak/net/msgstruct.hpp +++ b/xmrstak/net/msgstruct.hpp @@ -66,8 +66,9 @@ struct sock_err // Unlike socket errors, GPU errors are read-only strings struct gpu_res_err { + size_t idx; // GPU index const char* error_str; - gpu_res_err(const char* error_str) : error_str(error_str) {} + gpu_res_err(const char* error_str, size_t idx) : error_str(error_str), idx(idx) {} }; enum ex_event_name { EV_INVALID_VAL, EV_SOCK_READY, EV_SOCK_ERROR, EV_GPU_RES_ERROR, @@ -99,7 +100,7 @@ struct ex_event }; ex_event() { iName = EV_INVALID_VAL; iPoolId = 0;} - ex_event(const char* gpu_err, size_t id) : iName(EV_GPU_RES_ERROR), iPoolId(id), oGpuError(gpu_err) {} + ex_event(const char* gpu_err, size_t gpu_idx, size_t id) : iName(EV_GPU_RES_ERROR), iPoolId(id), oGpuError(gpu_err, gpu_idx) {} ex_event(std::string&& err, bool silent, size_t id) : iName(EV_SOCK_ERROR), iPoolId(id), oSocketError(std::move(err), silent) { } ex_event(job_result dat, size_t id) : iName(EV_MINER_HAVE_RESULT), iPoolId(id), oJobResult(dat) {} ex_event(pool_job dat, size_t id) : iName(EV_POOL_HAVE_JOB), iPoolId(id), oPoolJob(dat) {} |
