diff --git a/CMakeLists.txt b/CMakeLists.txt
index 67ab6ca..de75f39 100644
--- a/CMakeLists.txt
+++ b/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()"
+)
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-----
 ```
 
diff --git a/amd_gpu/gpu.c b/amd_gpu/gpu.c
index 7571e5d..a888aef 100644
--- a/amd_gpu/gpu.c
+++ b/amd_gpu/gpu.c
@@ -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);
diff --git a/cli-miner.cpp b/cli-miner.cpp
index 365ea27..252d888 100644
--- a/cli-miner.cpp
+++ b/cli-miner.cpp
@@ -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;
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<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);
 
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 <stdlib.h>
 #include <stdio.h>
 #include <string.h>
@@ -142,3 +144,5 @@ bool httpd::start_daemon()
 	return true;
 }
 
+#endif
+
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; }
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/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
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__)
 	
diff --git a/version.h b/version.h
new file mode 100644
index 0000000..57d5ff9
--- /dev/null
+++ b/version.h
@@ -0,0 +1,4 @@
+#pragma once
+
+#define XMR_STAK_NAME "xmr-stak-amd"
+#define XMR_STAK_VERSION "1.1.0-1.4.0"