11 files changed, 59 insertions, 10 deletions
diff --git a/.appveyor.yml b/.appveyor.yml
index 7989cb6..efa65bf 100644
--- a/.appveyor.yml
+++ b/.appveyor.yml
@@ -29,4 +29,4 @@ test_script:
   - cd c:\xmr-stak\build\bin\Release
   - dir
   - copy C:\xmr-stak-dep\openssl\bin\* .
-  - xmr-stak.exe --help
+#  - xmr-stak.exe --help
diff --git a/doc/compile.md b/doc/compile.md
index cf25135..4fa940f 100644
--- a/doc/compile.md
+++ b/doc/compile.md
@@ -9,6 +9,7 @@
 * [Compile on Windows](compile_Windows.md)
 * [Compile on Linux](compile_Linux.md)
 * [Compile on FreeBSD](compile_FreeBSD.md)
+* [Compile on MacOS](compile_MacOS.md)
 
 ## Build System
 
@@ -30,6 +31,7 @@ After the configuration you need to compile the miner, follow the guide for your
 * [Compile in Windows](compile_Windows.md)
 * [Compile in Linux](compile_Linux.md)
 * [Compile in FreeBSD](compile_FreeBSD.md)
+* [Compile in MacOS](compile_MacOS.md)
 
 ## Generic Build Options
 - `CMAKE_INSTALL_PREFIX` install miner to the home folder
diff --git a/doc/compile_MacOS.md b/doc/compile_MacOS.md
new file mode 100644
index 0000000..1b0af91
--- /dev/null
+++ b/doc/compile_MacOS.md
@@ -0,0 +1,31 @@
+# Compile **xmr-stak** for MacOS
+
+## Dependencies
+
+Assuming you already have [Homebrew](https://brew.sh) installed, the installation of dependencies is pretty straightforward and will generate the `xmr-stak` binary in the `bin/` directory.
+
+### For NVIDIA GPUs
+
+```shell
+brew tap caskroom/drivers
+brew cask install nvidia-cuda
+brew install hwloc libmicrohttpd gcc openssl cmake
+cmake . -DOPENSSL_ROOT_DIR=/usr/local/opt/openssl -DOpenCL_ENABLE=OFF
+make install
+```
+
+[All available CMake options](compile.md#nvidia-build-options)
+
+### For AMD GPUs
+
+> 🖐 We need help with AMD GPU compilation instructions. Please submit a PR if you managed to install [AMD APP SDK](http://developer.amd.com/amd-accelerated-parallel-processing-app-sdk/) and to compile `xmr-stak` on MacOS.
+
+### For CPU-only mining
+
+```shell
+brew install hwloc libmicrohttpd gcc openssl cmake
+cmake . -DOPENSSL_ROOT_DIR=/usr/local/opt/openssl -DCUDA_ENABLE=OFF -DOpenCL_ENABLE=OFF
+make install
+```
+
+[All available CMake options](compile.md#cpu-build-options)
diff --git a/doc/compile_Windows.md b/doc/compile_Windows.md
index 532348c..dd7af1e 100644
--- a/doc/compile_Windows.md
+++ b/doc/compile_Windows.md
@@ -81,15 +81,19 @@
   mkdir build
   cd build
   ```
-  - with CUDA 8
+  - for CUDA 8*
     ```
     cmake -G "Visual Studio 15 2017 Win64" -T v140,host=x64 ..
     ```
-  - with CUDA 9
+  - for CUDA 9 and/or AMD GPUs, CPU
     ```
     cmake -G "Visual Studio 15 2017 Win64" -T v141,host=x64 ..
     ```
   ```
   cmake --build . --config Release --target install
   cd bin\Release
+  copy C:\xmr-stak-dep\openssl\bin\* .
   ```
+
+\* Miner is also compiled for AMD GPUs (if the AMD APP SDK is installed) and CPUs.
+CUDA 8 requires a downgrade to the old v140 tool chain.
diff --git a/doc/usage.md b/doc/usage.md
index ca379ab..60825b4 100644
--- a/doc/usage.md
+++ b/doc/usage.md
@@ -20,7 +20,7 @@ The number of files depends on the available backends.
 1) Double click the `xmr-stak.exe` file
 2) Fill in the pool url, username and password
 
-## Usage on Linux
+## Usage on Linux & MacOS
 1) Open a terminal within the folder with the binary
 2) Start the miner with `./xmr-stak`
 
@@ -40,6 +40,8 @@ Usage: xmr-stak [OPTION]...
   --cpu FILE            CPU backend miner config file
   --noAMD               disable the AMD miner backend
   --amd FILE            AMD backend miner config file
+  --noNVIDIA            disable the NVIDIA miner backend
+  --nvidia FILE         NVIDIA backend miner config file
 
 The Following options temporary overwrites the config file settings:
   -o, --url URL         pool url and port, e.g. pool.usxmrpool.com:3333
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp
index 15b8457..879a2e4 100644
--- a/xmrstak/backend/amd/amd_gpu/gpu.cpp
+++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp
@@ -476,7 +476,7 @@ std::vector<GpuContext> getAMDDevices(int index)
 						if(clStatus == CL_SUCCESS)
 						{
 							std::string devVendor(devVendorVec.data());
-							if( devVendor.find("Advanced Micro Devices") != std::string::npos)
+							if( devVendor.find("Advanced Micro Devices") != std::string::npos || devVendor.find("AMD") != std::string::npos)
 							{
 								GpuContext ctx;
 								ctx.deviceIdx = k;
@@ -541,7 +541,7 @@ 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)
+			if( platformName.find("Advanced Micro Devices") != std::string::npos || platformName.find("Apple") != std::string::npos)
 			{
 				platformIndex = i;
 				printer::inst()->print_msg(L0,"Found AMD platform index id = %i, name = %s",i , platformName.c_str());
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.hpp b/xmrstak/backend/amd/amd_gpu/gpu.hpp
index 123de01..abbd08d 100644
--- a/xmrstak/backend/amd/amd_gpu/gpu.hpp
+++ b/xmrstak/backend/amd/amd_gpu/gpu.hpp
@@ -9,6 +9,7 @@
 #endif
 
 #include <stdint.h>
+#include <string>
 #include <vector>
 
 #define ERR_SUCCESS (0)
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
index a92fa8c..dba6676 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
@@ -167,10 +167,10 @@ __forceinline__ __device__ uint32_t shuffle(volatile uint32_t* ptr,const uint32_
 #endif
 }
 
+template<size_t ITERATIONS, uint32_t THREAD_SHIFT, uint32_t MASK>
 #ifdef XMR_STAK_THREADS
 __launch_bounds__( XMR_STAK_THREADS * 4 )
 #endif
-template<size_t ITERATIONS, uint32_t THREAD_SHIFT, uint32_t MASK>
 __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int partidx, uint32_t * d_long_state, uint32_t * d_ctx_a, uint32_t * d_ctx_b )
 {
 	__shared__ uint32_t sharedMemory[1024];
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
index 5501d8d..333ae73 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
@@ -380,6 +380,10 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
 		 */
 		ctx->device_blocks = props.multiProcessorCount *
 			( props.major < 3 ? 2 : 3 );
+
+		// increase bfactor for low end devices to avoid that the miner is killed by the OS
+		if(props.multiProcessorCount < 6)
+			ctx->device_bfactor += 2;
 	}
 	if(ctx->device_threads == -1)
 	{
diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp
index 462f564..34bde6c 100644
--- a/xmrstak/jconf.cpp
+++ b/xmrstak/jconf.cpp
@@ -158,6 +158,7 @@ bool jconf::GetPoolConfig(size_t id, pool_cfg& cfg)
 	}
 	else /* Special case - user selected same weights for everything */
 		cfg.weight = 0.0;
+	return true;
 }
 
 bool jconf::TlsSecureAlgos()
diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp
index cccfca7..ca78c61 100644
--- a/xmrstak/misc/executor.cpp
+++ b/xmrstak/misc/executor.cpp
@@ -459,6 +459,7 @@ void executor::on_miner_result(size_t pool_id, job_result& oResult)
 void disable_sigpipe()
 {
 	struct sigaction sa;
+	memset(&sa, 0, sizeof(sa));
 	sa.sa_handler = SIG_IGN;
 	sa.sa_flags = 0;
 	if (sigaction(SIGPIPE, &sa, 0) == -1)
@@ -624,7 +625,10 @@ inline const char* hps_format(double h, char* buf, size_t l)
 {
 	if(std::isnormal(h) || h == 0.0)
 	{
-		snprintf(buf, l, " %03.1f", h);
+		if(h < 10.0)
+			snprintf(buf, l, "  %03.1f", h);
+		else
+			snprintf(buf, l, " %04.1f", h);
 		return buf;
 	}
 	else
@@ -722,9 +726,9 @@ void executor::hashrate_report(std::string& out)
 			std::transform(name.begin(), name.end(), name.begin(), ::toupper);
 			
 			out.append("HASHRATE REPORT - ").append(name).append("\n");
-			out.append("| ID | 10s |  60s |  15m |");
+			out.append("| ID |  10s |  60s |  15m |");
 			if(nthd != 1)
-				out.append(" ID | 10s |  60s |  15m |\n");
+				out.append(" ID |  10s |  60s |  15m |\n");
 			else
 				out.append(1, '\n');