Merge pull request #1720 from fireice-uk/dev

release 2.4.6

README.md

@@ -1,9 +1,7 @@
 ###### fireice-uk's and psychocrypt's
-# XMR-Stak - Monero/Aeon All-in-One Mining Software
+# XMR-Stak - Cryptonight All-in-One Mining Software
 
-**XMR-Stak is ready for the POW change of Monero-v7, Aeon-v7, stellite-v4 and Sumukoin-v3**
-
-XMR-Stak is a universal Stratum pool miner. This miner supports CPUs, AMD and NVIDIA gpus and can be used to mine the crypto currency Monero and Aeon.
+XMR-Stak is a universal Stratum pool miner. This miner supports CPUs, AMD and NVIDIA gpus and can be used to mine the crypto currencys Monero, Aeon and many more Cryptonight coins.
 
 ## HTML reports
 <img src="https://gist.githubusercontent.com/fireice-uk/2da301131ac01695ff79539a27b81d68/raw/4c09cdeee86f94df2e9dd86b927e64aded6184f5/xmr-stak-cpu-hashrate.png" width="260"> <img src="https://gist.githubusercontent.com/fireice-uk/2da301131ac01695ff79539a27b81d68/raw/4c09cdeee86f94df2e9dd86b927e64aded6184f5/xmr-stak-cpu-results.png" width="260"> <img src="https://gist.githubusercontent.com/fireice-uk/2da301131ac01695ff79539a27b81d68/raw/4c09cdeee86f94df2e9dd86b927e64aded6184f5/xmr-stak-cpu-connection.png" width="260">
@@ -42,22 +40,28 @@ Besides [Monero](https://getmonero.org), following coins can be mined using this
 
 - [Aeon](http://www.aeon.cash)
 - [BBSCoin](https://www.bbscoin.xyz)
-- [Croat](https://croat.cat)
-- [Edollar](https://edollar.cash)
-- [Electroneum](https://electroneum.com)
+- [BitTube](https://coin.bit.tube/)
 - [Graft](https://www.graft.network)
 - [Haven](https://havenprotocol.com)
 - [Intense](https://intensecoin.com)
-- [IPBC](https://ipbc.io)
-- [Karbo](https://karbo.io)
 - [Masari](https://getmasari.org)
-- [Sumokoin](https://www.sumokoin.org)
+- [Ryo](https://ryo-currency.com)
 - [TurtleCoin](https://turtlecoin.lol)
 
 If your prefered coin is not listed, you can chose one of the following algorithms:
 
-- Cryptonight - 2 MiB scratchpad memory
-- Cryptonight-light - 1 MiB scratchpad memory
+- 1MiB scratchpad memory
+    - cryptonight_lite
+    - cryptonight_lite_v7
+    - cryptonight_lite_v7_xor (algorithm used by ipbc)
+- 2MiB scratchpad memory
+    - cryptonight
+    - cryptonight_masari
+    - cryptonight_v7
+    - cryptonight_v7_stellite
+- 4MiB scratchpad memory
+    - cryptonight_haven
+    - cryptonight_heavy
 
 Please note, this list is not complete, and is not an endorsement.
 

doc/FAQ.md

@@ -35,7 +35,7 @@ Reference: http://rybkaforum.net/cgi-bin/rybkaforum/topic_show.pl?pid=259791#pid
 
 If you set up the user rights properly ([see above](https://github.com/fireice-uk/xmr-stak/blob/master/doc/FAQ.md#selockmemoryprivilege-failed)), and your system has 4-8GB of RAM (50%+ use), there is a significant chance that there simply won't be a large enough chunk of contiguous memory because Windows is fairly bad at mitigating memory fragmentation.
 
-If that happens, disable all auto-staring applications and run the miner after a reboot.
+If that happens, disable all auto-starting applications and run the miner after a reboot.
 
 ## Error msvcp140.dll and vcruntime140.dll not available
 
@@ -46,11 +46,11 @@ Download and install this [runtime package](https://go.microsoft.com/fwlink/?Lin
 
 On Linux you will need to configure large page support and increase your ulimit -l.
 
-To set large page support, add the following lines to /etc/sysctl.conf:
+To set large page support, add the following lines to `/etc/sysctl.conf` (`/etc/sysctl.d/xmr-stak.conf` for [Arch Linux](https://www.archlinux.org/news/deprecation-of-etcsysctlconf/) and its derivatives):
 
     vm.nr_hugepages=128
 
-To increase the ulimit, add following lines to /etc/security/limits.conf:
+To increase the ulimit, add following lines to `/etc/security/limits.conf`:
 
     * soft memlock 262144
     * hard memlock 262144

doc/compile.md

@@ -23,8 +23,8 @@ There are two easy ways to set variables for `cmake` to configure *xmr-stak*
   - edit your options
   - end the GUI by pressing the key `c`(create) and than `g`(generate)
 - set Options on the command line
-  - enable a option: `cmake .. -DNAME_OF_THE_OPTION=ON`
-  - disable a option `cmake .. -DNAME_OF_THE_OPTION=OFF`
+  - enable an option: `cmake .. -DNAME_OF_THE_OPTION=ON`
+  - disable an option `cmake .. -DNAME_OF_THE_OPTION=OFF`
   - set a value `cmake .. -DNAME_OF_THE_OPTION=value`
 
 After the configuration you need to compile the miner, follow the guide for your platform:
@@ -43,28 +43,28 @@ After the configuration you need to compile the miner, follow the guide for your
 - `CMAKE_BUILD_TYPE` set the build type
   - valid options: `Release` or `Debug`
   - you should always keep `Release` for your productive miners
-- `MICROHTTPD_ENABLE` allow to disable/enable the dependency *microhttpd*
+- `MICROHTTPD_ENABLE` allows to disable/enable the dependency *microhttpd*
   - there is no *http* interface available if option is disabled: `cmake .. -DMICROHTTPD_ENABLE=OFF`
-- `OpenSSL_ENABLE` allow to disable/enable the dependency *OpenSSL*
+- `OpenSSL_ENABLE` allows to disable/enable the dependency *OpenSSL*
   - it is not possible to connect to a *https* secured pool if option is disabled: `cmake .. -DOpenSSL_ENABLE=OFF`
 - `XMR-STAK_COMPILE` select the CPU compute architecture (default: native)
   - native means the miner binary can be used only on the system where it is compiled but will archive the highest hash rate
   - use `cmake .. -DXMR-STAK_COMPILE=generic` to run the miner on all CPU's with sse2
 
 ## CPU Build Options
 
-- `CPU_ENABLE` allow to disable/enable the CPU backend of the miner
-- `HWLOC_ENABLE` allow to disable/enable the dependency *hwloc*
+- `CPU_ENABLE` allows to disable/enable the CPU backend of the miner
+- `HWLOC_ENABLE` allows to disable/enable the dependency *hwloc*
   - the config suggestion is not optimal if option is disabled: `cmake .. -DHWLOC_ENABLE=OFF`
   - disabling can be reduce the miner performance
 
 ## AMD Build Options
 
-- `OpenCL_ENABLE` allow to disable/enable the AMD backend of the miner
+- `OpenCL_ENABLE` allows to disable/enable the AMD backend of the miner
 
 ## NVIDIA Build Options
 
-- `CUDA_ENABLE` allow to disable/enable the NVIDIA backend of the miner
+- `CUDA_ENABLE` allows to disable/enable the NVIDIA backend of the miner
 - `CUDA_ARCH` build for a certain compute architecture
   - this option needs a semicolon separated list
   - `cmake .. -DCUDA_ARCH=61` or `cmake .. -DCUDA_ARCH=20;61`

scripts/build_xmr-stak_docker/build_xmr-stak_docker.sh

@@ -18,7 +18,7 @@ chmod a+x cuda_*_linux-run
 ########################
 # Fedora 27
 ########################
-# CUDA is not going to work on Fedora 27 beacuse it's only support these distributions: http://docs.nvidia.com/cuda/cuda-installation-guide-linux/index.html
+# CUDA is not going to work on Fedora 27 beacuse it only supports these distributions: http://docs.nvidia.com/cuda/cuda-installation-guide-linux/index.html
 docker run --rm -it -v $PWD:/mnt fedora:27 /bin/bash -c "
 set -x ;
 dnf install -y -q cmake gcc-c++ hwloc-devel libmicrohttpd-devel libstdc++-static make openssl-devel;

xmrstak/backend/amd/amd_gpu/gpu.cpp

@@ -1004,7 +1004,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
 		return(ERR_OCL_API);
 	}
 
-	if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon || miner_algo == cryptonight_ipbc || miner_algo == cryptonight_stellite || miner_algo == cryptonight_masari)
+	if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon || miner_algo == cryptonight_ipbc || miner_algo == cryptonight_stellite || miner_algo == cryptonight_masari || miner_algo == cryptonight_bittube2)
 	{
 		// Input
 		if ((ret = clSetKernelArg(ctx->Kernels[kernel_storage][1], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS)

xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl

@@ -513,16 +513,17 @@ __kernel void JOIN(cn0,ALGO)(__global ulong *input, __global uint4 *Scratchpad,
 
 	mem_fence(CLK_LOCAL_MEM_FENCE);
 
-// cryptonight_heavy or cryptonight_haven
-#if (ALGO == 4 || ALGO == 9)
+// cryptonight_heavy || cryptonight_haven || cryptonight_bittube2
+#if (ALGO == 4 || ALGO == 9 || ALGO == 10)
 	__local uint4 xin[8][WORKSIZE];
 
 	/* Also left over threads perform this loop.
 	 * The left over thread results will be ignored
 	 */
+	#pragma unroll 16
 	for(size_t i=0; i < 16; i++)
 	{
-		#pragma unroll
+		#pragma unroll 10
 		for(int j = 0; j < 10; ++j)
 			text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey1)[j]);
 		barrier(CLK_LOCAL_MEM_FENCE);
@@ -550,11 +551,11 @@ __kernel void JOIN(cn0,ALGO)(__global ulong *input, __global uint4 *Scratchpad,
 	}
 	mem_fence(CLK_GLOBAL_MEM_FENCE);
 }
-
+		
 __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
 __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states, ulong Threads
-// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_masari
-#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8)
+// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_masari || cryptonight_bittube2
+#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8 || ALGO == 10)
 , __global ulong *input
 #endif
 )
@@ -574,8 +575,8 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
 	}
 
 	barrier(CLK_LOCAL_MEM_FENCE);
-// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_masari
-#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8)
+// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_masari || cryptonight_bittube2
+#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8 || ALGO == 10)
     uint2 tweak1_2;
 #endif
 	uint4 b_x;
@@ -599,8 +600,8 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
 		b[1] = states[3] ^ states[7];
 
 		b_x = ((uint4 *)b)[0];
-// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_masari
-#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8)
+// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_masari || cryptonight_bittube2
+#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8 || ALGO == 10)
 		tweak1_2 = as_uint2(input[4]);
 		tweak1_2.s0 >>= 24;
 		tweak1_2.s0 |= tweak1_2.s1 << 8;
@@ -624,11 +625,15 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
 			ulong c[2];
 
 			((uint4 *)c)[0] = Scratchpad[IDX((idx0 & MASK) >> 4)];
+// cryptonight_bittube2
+#if(ALGO == 10)
+			((uint4 *)c)[0] = AES_Round_bittube2(AES0, AES1, AES2, AES3, ((uint4 *)c)[0], ((uint4 *)a)[0]);
+#else
 			((uint4 *)c)[0] = AES_Round(AES0, AES1, AES2, AES3, ((uint4 *)c)[0], ((uint4 *)a)[0]);
-
+#endif
 			b_x ^= ((uint4 *)c)[0];
-// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_masari
-#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8)
+// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_masari || cryptonight_bittube2
+#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8 || ALGO == 10)
 			uint table = 0x75310U;
 // cryptonight_stellite
 #	if(ALGO == 7)
@@ -646,10 +651,11 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
 			a[1] += c[0] * as_ulong2(tmp).s0;
 			a[0] += mul_hi(c[0], as_ulong2(tmp).s0);
 
-// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_masari
-#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8)
+// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_masari || cryptonight_bittube2
+#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8 || ALGO == 10)
 
-#	if(ALGO == 6)
+// cryptonight_ipbc || cryptonight_bittube2
+#	if(ALGO == 6 || ALGO == 10)
 			uint2 ipbc_tmp = tweak1_2 ^ ((uint2 *)&(a[0]))[0];
 			((uint2 *)&(a[1]))[0] ^= ipbc_tmp;
 			Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0];
@@ -669,8 +675,8 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
 
 			b_x = ((uint4 *)c)[0];
 
-// cryptonight_heavy
-#if (ALGO == 4)
+// cryptonight_heavy || cryptonight_bittube2
+#if (ALGO == 4 || ALGO == 10)
 			long n = *((__global long*)(Scratchpad + (IDX((idx0 & MASK) >> 4))));
 			int d = ((__global int*)(Scratchpad + (IDX((idx0 & MASK) >> 4))))[2];
 			long q = n / (d | 0x5);
@@ -743,8 +749,8 @@ __kernel void JOIN(cn2,ALGO) (__global uint4 *Scratchpad, __global ulong *states
 	}
 
 	barrier(CLK_LOCAL_MEM_FENCE);
-// cryptonight_heavy or cryptonight_haven
-#if (ALGO == 4 || ALGO == 9)
+// cryptonight_heavy || cryptonight_haven || cryptonight_bittube2
+#if (ALGO == 4 || ALGO == 9 || ALGO == 10)
 	__local uint4 xin[8][WORKSIZE];
 #endif
 
@@ -753,8 +759,8 @@ __kernel void JOIN(cn2,ALGO) (__global uint4 *Scratchpad, __global ulong *states
 	if(gIdx < Threads)
 #endif
 	{
-// cryptonight_heavy or cryptonight_haven
-#if (ALGO == 4 || ALGO == 9)
+// cryptonight_heavy || cryptonight_haven || cryptonight_bittube2
+#if (ALGO == 4 || ALGO == 9 || ALGO == 10)
 		#pragma unroll 2
 		for(int i = 0; i < (MEMORY >> 7); ++i)
 		{
@@ -800,14 +806,15 @@ __kernel void JOIN(cn2,ALGO) (__global uint4 *Scratchpad, __global ulong *states
 #endif
 	}
 
-// cryptonight_heavy or cryptonight_haven
-#if (ALGO == 4 || ALGO == 9)
+// cryptonight_heavy or cryptonight_haven || cryptonight_bittube2
+#if (ALGO == 4 || ALGO == 9 || ALGO == 10)
 	/* Also left over threads perform this loop.
 	 * The left over thread results will be ignored
 	 */
+	#pragma unroll 16
 	for(size_t i=0; i < 16; i++)
 	{
-		#pragma unroll
+		#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);

xmrstak/backend/amd/amd_gpu/opencl/wolf-aes.cl

@@ -74,6 +74,19 @@ static const __constant uint AES0_C[256] =
 
 #define BYTE(x, y)	(amd_bfe((x), (y) << 3U, 8U))
 
+inline uint4 AES_Round_bittube2(const __local uint *AES0, const __local uint *AES1, const __local uint *AES2, const __local uint *AES3, uint4 x, uint4 k)
+{
+	x = ~x;
+	k[0] ^= AES0[BYTE(x[0], 0)] ^ AES1[BYTE(x[1], 1)] ^ AES2[BYTE(x[2], 2)] ^ AES3[BYTE(x[3], 3)];
+	x[0] ^= k[0];
+	k[1] ^= AES0[BYTE(x[1], 0)] ^ AES1[BYTE(x[2], 1)] ^ AES2[BYTE(x[3], 2)] ^ AES3[BYTE(x[0], 3)];
+	x[1] ^= k[1];
+	k[2] ^= AES0[BYTE(x[2], 0)] ^ AES1[BYTE(x[3], 1)] ^ AES2[BYTE(x[0], 2)] ^ AES3[BYTE(x[1], 3)];
+	x[2] ^= k[2];
+	k[3] ^= AES0[BYTE(x[3], 0)] ^ AES1[BYTE(x[0], 1)] ^ AES2[BYTE(x[1], 2)] ^ AES3[BYTE(x[2], 3)];
+	return k;
+}
+
 uint4 AES_Round(const __local uint *AES0, const __local uint *AES1, const __local uint *AES2, const __local uint *AES3, const uint4 X, uint4 key)
 {
 	key.s0 ^= AES0[BYTE(X.s0, 0)];

xmrstak/backend/amd/minethd.cpp

@@ -254,7 +254,7 @@ void minethd::work_main()
 
 				hash_fun(bWorkBlob, oWork.iWorkSize, bResult, cpu_ctx);
 				if ( (*((uint64_t*)(bResult + 24))) < oWork.iTarget)
-					executor::inst()->push_event(ex_event(job_result(oWork.sJobID, results[i], bResult, iThreadNo), oWork.iPoolId));
+					executor::inst()->push_event(ex_event(job_result(oWork.sJobID, results[i], bResult, iThreadNo, miner_algo), oWork.iPoolId));
 				else
 					executor::inst()->push_event(ex_event("AMD Invalid Result", pGpuCtx->deviceIdx, oWork.iPoolId));
 			}