Merge pull request #1459 from fireice-uk/dev2.4.3

release 2.4.3

54 files changed, 667 insertions, 613 deletions

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 15a2684..784c0bd 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -148,6 +148,11 @@ if(CUDA_ENABLE)
                 set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -std=c++11")
             endif()
 
+            # avoid that nvcc in CUDA 8 complains about sm_20 pending removal
+	    if(CUDA_VERSION VERSION_EQUAL 8.0)
+                set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -Wno-deprecated-gpu-targets")
+            endif()
+
             # avoid that nvcc in CUDA < 8 tries to use libc `memcpy` within the kernel
             if(CUDA_VERSION VERSION_LESS 8.0)
                 add_definitions(-D_FORCE_INLINES)
@@ -522,8 +527,8 @@ else()
     add_executable(xmr-stak ${SRCFILES_CPP})
 endif()
 
-set(EXECUTABLE_OUTPUT_PATH "bin")
-set(LIBRARY_OUTPUT_PATH "bin")
+set(EXECUTABLE_OUTPUT_PATH "bin" CACHE STRING "Path to place executables relative to ${CMAKE_INSTALL_PREFIX}")
+set(LIBRARY_OUTPUT_PATH "bin" CACHE STRING "Path to place libraries relative to ${CMAKE_INSTALL_PREFIX}")
 
 target_link_libraries(xmr-stak ${LIBS} xmr-stak-c xmr-stak-backend)
 
@@ -535,23 +540,23 @@ target_link_libraries(xmr-stak ${LIBS} xmr-stak-c xmr-stak-backend)
 # 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
-            RUNTIME DESTINATION "${CMAKE_INSTALL_PREFIX}/bin")
+            RUNTIME DESTINATION "${CMAKE_INSTALL_PREFIX}/${EXECUTABLE_OUTPUT_PATH}")
     if(CUDA_FOUND)
         if(WIN32)
             install(TARGETS xmrstak_cuda_backend
-                RUNTIME DESTINATION "${CMAKE_INSTALL_PREFIX}/bin")
+                RUNTIME DESTINATION "${CMAKE_INSTALL_PREFIX}/${LIBRARY_OUTPUT_PATH}")
         else()
             install(TARGETS xmrstak_cuda_backend
-                LIBRARY DESTINATION "${CMAKE_INSTALL_PREFIX}/bin")
+                LIBRARY DESTINATION "${CMAKE_INSTALL_PREFIX}/${LIBRARY_OUTPUT_PATH}")
         endif()
     endif()
     if(OpenCL_FOUND)
         if(WIN32)
             install(TARGETS xmrstak_opencl_backend
-                RUNTIME DESTINATION "${CMAKE_INSTALL_PREFIX}/bin")
+                RUNTIME DESTINATION "${CMAKE_INSTALL_PREFIX}/${LIBRARY_OUTPUT_PATH}")
         else()
             install(TARGETS xmrstak_opencl_backend
-                LIBRARY DESTINATION "${CMAKE_INSTALL_PREFIX}/bin")
+                LIBRARY DESTINATION "${CMAKE_INSTALL_PREFIX}/${LIBRARY_OUTPUT_PATH}")
         endif()
     endif()
 else()
diff --git a/README.md b/README.md
index 047aa51..64fd488 100644
--- a/README.md
+++ b/README.md
@@ -41,9 +41,12 @@ XMR-Stak is a universal Stratum pool miner. This miner supports CPUs, AMD and NV
 Besides [Monero](https://getmonero.org), following coins can be mined using this miner:
 
 - [Aeon](http://www.aeon.cash)
+- [BBSCoin](https://www.bbscoin.xyz)
+- [Croat](https://croat.cat)
 - [Edollar](https://edollar.cash)
 - [Electroneum](https://electroneum.com)
 - [Graft](https://www.graft.network)
+- [Haven](https://havenprotocol.com)
 - [Intense](https://intensecoin.com)
 - [Karbo](https://karbo.io)
 - [Sumokoin](https://www.sumokoin.org)
diff --git a/doc/FAQ.md b/doc/FAQ.md
index 35dc248..0661073 100644
--- a/doc/FAQ.md
+++ b/doc/FAQ.md
@@ -9,7 +9,7 @@
 * [Virus Protection Alert](#virus-protection-alert)
 * [Change Currency to Mine](#change-currency-to-mine)
 * [How can I mine Monero](#how-can-i-mine-monero)
-* [Why is Monero named monero7](why-is-monero-named-monero7)
+* [Why is Monero named monero7](#why-is-monero-named-monero7)
 * [Which currency must be chosen if my fork coin is not listed](#which-currency-must-be-chosen-if-my-fork-coin-is-not-listed)
 
 ## "Obtaining SeLockMemoryPrivilege failed."
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp
index 006a7ed..03100d0 100644
--- a/xmrstak/backend/amd/amd_gpu/gpu.cpp
+++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp
@@ -26,6 +26,7 @@
 #include <algorithm>
 #include <regex>
 #include <cassert>
+#include <algorithm> 
 
 #include <fstream>
 #include <sstream>
@@ -307,12 +308,13 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_
 		return ERR_OCL_API;
 	}
 
-	size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
-	int threadMemMask = cn_select_mask(::jconf::inst()->GetMiningAlgo());
-	int hashIterations = cn_select_iter(::jconf::inst()->GetMiningAlgo());
+	size_t scratchPadSize = std::max(
+		cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()),
+		cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot())
+	);
 
 	size_t g_thd = ctx->rawIntensity;
-	ctx->ExtraBuffers[0] = clCreateBuffer(opencl_ctx, CL_MEM_READ_WRITE, hashMemSize * g_thd, NULL, &ret);
+	ctx->ExtraBuffers[0] = clCreateBuffer(opencl_ctx, CL_MEM_READ_WRITE, scratchPadSize * g_thd, NULL, &ret);
 	if(ret != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clCreateBuffer to create hash scratchpads buffer.", err_to_str(ret));
@@ -373,167 +375,203 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_
 		return ERR_OCL_API;
 	}
 
-	auto miner_algo = ::jconf::inst()->GetMiningAlgo();
-
-	char options[512];
-	snprintf(options, sizeof(options),
-		"-DITERATIONS=%d -DMASK=%d -DWORKSIZE=%llu -DSTRIDED_INDEX=%d -DMEM_CHUNK_EXPONENT=%d  -DCOMP_MODE=%d -DMEMORY=%llu -DALGO=%d",
-		hashIterations, threadMemMask, int_port(ctx->workSize), ctx->stridedIndex, int(1u<<ctx->memChunk), ctx->compMode ? 1 : 0, 
-		int_port(hashMemSize), int(miner_algo));
-	/* create a hash for the compile time cache
-	 * used data:
-	 *   - source code
-	 *   - device name
-	 *   - compile paramater
-	 */
-	std::string src_str(source_code);
-	src_str += options;
-	src_str += devNameVec.data();
-	std::string hash_hex_str;
-	picosha2::hash256_hex_string(src_str, hash_hex_str);
-
-	std::string cache_file = get_home() + "/.openclcache/" + hash_hex_str + ".openclbin";
-	std::ifstream clBinFile(cache_file, std::ofstream::in | std::ofstream::binary);
-	if(xmrstak::params::inst().AMDCache == false || !clBinFile.good())
-	{
-		if(xmrstak::params::inst().AMDCache)
-			printer::inst()->print_msg(L1,"OpenCL device %u - Precompiled code %s not found. Compiling ...",ctx->deviceIdx, cache_file.c_str());
-		ctx->Program = clCreateProgramWithSource(opencl_ctx, 1, (const char**)&source_code, NULL, &ret);
-		if(ret != CL_SUCCESS)
-		{
-			printer::inst()->print_msg(L1,"Error %s when calling clCreateProgramWithSource on the OpenCL miner code", err_to_str(ret));
-			return ERR_OCL_API;
-		}
-
-		ret = clBuildProgram(ctx->Program, 1, &ctx->DeviceID, options, NULL, NULL);
-		if(ret != CL_SUCCESS)
+	xmrstak_algo miner_algo[2] = {
+		::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo(),
+		::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot()
+	};
+	int num_algos = miner_algo[0] == miner_algo[1] ? 1 : 2;
+
+	for(int ii = 0; ii < num_algos; ++ii)
+	{
+		// scratchpad size for the selected mining algorithm
+		size_t hashMemSize = cn_select_memory(miner_algo[ii]);
+		int threadMemMask = cn_select_mask(miner_algo[ii]);
+		int hashIterations = cn_select_iter(miner_algo[ii]);
+
+		char options[512];
+		snprintf(options, sizeof(options),
+			"-DITERATIONS=%d -DMASK=%d -DWORKSIZE=%llu -DSTRIDED_INDEX=%d -DMEM_CHUNK_EXPONENT=%d  -DCOMP_MODE=%d -DMEMORY=%llu -DALGO=%d",
+		hashIterations, threadMemMask, int_port(ctx->workSize), ctx->stridedIndex, int(1u<<ctx->memChunk), ctx->compMode ? 1 : 0,
+			int_port(hashMemSize), int(miner_algo[ii]));
+		/* create a hash for the compile time cache
+		 * used data:
+		 *   - source code
+		 *   - device name
+		 *   - compile paramater
+		 */
+		std::string src_str(source_code);
+		src_str += options;
+		src_str += devNameVec.data();
+		std::string hash_hex_str;
+		picosha2::hash256_hex_string(src_str, hash_hex_str);
+
+		std::string cache_file = get_home() + "/.openclcache/" + hash_hex_str + ".openclbin";
+		std::ifstream clBinFile(cache_file, std::ofstream::in | std::ofstream::binary);
+		if(xmrstak::params::inst().AMDCache == false || !clBinFile.good())
 		{
-			size_t len;
-			printer::inst()->print_msg(L1,"Error %s when calling clBuildProgram.", err_to_str(ret));
-
-			if((ret = clGetProgramBuildInfo(ctx->Program, ctx->DeviceID, CL_PROGRAM_BUILD_LOG, 0, NULL, &len)) != CL_SUCCESS)
+			if(xmrstak::params::inst().AMDCache)
+				printer::inst()->print_msg(L1,"OpenCL device %u - Precompiled code %s not found. Compiling ...",ctx->deviceIdx, cache_file.c_str());
+			ctx->Program[ii] = clCreateProgramWithSource(opencl_ctx, 1, (const char**)&source_code, NULL, &ret);
+			if(ret != CL_SUCCESS)
 			{
-				printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for length of build log output.", err_to_str(ret));
+				printer::inst()->print_msg(L1,"Error %s when calling clCreateProgramWithSource on the OpenCL miner code", err_to_str(ret));
 				return ERR_OCL_API;
 			}
 
-			char* BuildLog = (char*)malloc(len + 1);
-			BuildLog[0] = '\0';
-
-			if((ret = clGetProgramBuildInfo(ctx->Program, ctx->DeviceID, CL_PROGRAM_BUILD_LOG, len, BuildLog, NULL)) != CL_SUCCESS)
+			ret = clBuildProgram(ctx->Program[ii], 1, &ctx->DeviceID, options, NULL, NULL);
+			if(ret != CL_SUCCESS)
 			{
+				size_t len;
+				printer::inst()->print_msg(L1,"Error %s when calling clBuildProgram.", err_to_str(ret));
+
+				if((ret = clGetProgramBuildInfo(ctx->Program[ii], ctx->DeviceID, CL_PROGRAM_BUILD_LOG, 0, NULL, &len)) != CL_SUCCESS)
+				{
+					printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for length of build log output.", err_to_str(ret));
+					return ERR_OCL_API;
+				}
+
+				char* BuildLog = (char*)malloc(len + 1);
+				BuildLog[0] = '\0';
+
+				if((ret = clGetProgramBuildInfo(ctx->Program[ii], ctx->DeviceID, CL_PROGRAM_BUILD_LOG, len, BuildLog, NULL)) != CL_SUCCESS)
+				{
+					free(BuildLog);
+					printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for build log.", err_to_str(ret));
+					return ERR_OCL_API;
+				}
+
+				printer::inst()->print_str("Build log:\n");
+				std::cerr<<BuildLog<<std::endl;
+
 				free(BuildLog);
-				printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for build log.", err_to_str(ret));
 				return ERR_OCL_API;
 			}
 
-			printer::inst()->print_str("Build log:\n");
-			std::cerr<<BuildLog<<std::endl;
+			cl_uint num_devices;
+			clGetProgramInfo(ctx->Program[ii], CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint), &num_devices,NULL);
 
-			free(BuildLog);
-			return ERR_OCL_API;
-		}
 
-		cl_uint num_devices;
-		clGetProgramInfo(ctx->Program, CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint), &num_devices,NULL);
-
-
-		std::vector<cl_device_id> devices_ids(num_devices);
-		clGetProgramInfo(ctx->Program, CL_PROGRAM_DEVICES, sizeof(cl_device_id)* devices_ids.size(), devices_ids.data(),NULL);
-		int dev_id = 0;
-		/* Search for the gpu within the program context.
-		 * The id can be different to  ctx->DeviceID.
-		 */
-		for(auto & ocl_device : devices_ids)
-		{
-			if(ocl_device == ctx->DeviceID)
-				break;
-			dev_id++;
-		}
+			std::vector<cl_device_id> devices_ids(num_devices);
+			clGetProgramInfo(ctx->Program[ii], CL_PROGRAM_DEVICES, sizeof(cl_device_id)* devices_ids.size(), devices_ids.data(),NULL);
+			int dev_id = 0;
+			/* Search for the gpu within the program context.
+			 * The id can be different to  ctx->DeviceID.
+			 */
+			for(auto & ocl_device : devices_ids)
+			{
+				if(ocl_device == ctx->DeviceID)
+					break;
+				dev_id++;
+			}
 
-		cl_build_status status;
-		do
-		{
-			if((ret = clGetProgramBuildInfo(ctx->Program, ctx->DeviceID, CL_PROGRAM_BUILD_STATUS, sizeof(cl_build_status), &status, NULL)) != CL_SUCCESS)
+			cl_build_status status;
+			do
 			{
-				printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for status of build.", err_to_str(ret));
-				return ERR_OCL_API;
+				if((ret = clGetProgramBuildInfo(ctx->Program[ii], ctx->DeviceID, CL_PROGRAM_BUILD_STATUS, sizeof(cl_build_status), &status, NULL)) != CL_SUCCESS)
+				{
+					printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for status of build.", err_to_str(ret));
+					return ERR_OCL_API;
+				}
+				port_sleep(1);
 			}
-			port_sleep(1);
-		}
-		while(status == CL_BUILD_IN_PROGRESS);
+			while(status == CL_BUILD_IN_PROGRESS);
 
-		std::vector<size_t> binary_sizes(num_devices);
-		clGetProgramInfo (ctx->Program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t) * binary_sizes.size(), binary_sizes.data(), NULL);
+			std::vector<size_t> binary_sizes(num_devices);
+			clGetProgramInfo (ctx->Program[ii], CL_PROGRAM_BINARY_SIZES, sizeof(size_t) * binary_sizes.size(), binary_sizes.data(), NULL);
 
-		std::vector<char*> all_programs(num_devices);
-		std::vector<std::vector<char>> program_storage;
+			std::vector<char*> all_programs(num_devices);
+			std::vector<std::vector<char>> program_storage;
+
+			if(xmrstak::params::inst().AMDCache)
+			{
+				int p_id = 0;
+				size_t mem_size = 0;
+				// create memory  structure to query all OpenCL program binaries
+				for(auto & p : all_programs)
+				{
+					program_storage.emplace_back(std::vector<char>(binary_sizes[p_id]));
+					all_programs[p_id] = program_storage[p_id].data();
+					mem_size += binary_sizes[p_id];
+					p_id++;
+				}
+
+				if((ret = clGetProgramInfo(ctx->Program[ii], CL_PROGRAM_BINARIES, num_devices * sizeof(char*), all_programs.data(),NULL)) != CL_SUCCESS)
+				{
+					printer::inst()->print_msg(L1,"Error %s when calling clGetProgramInfo.", err_to_str(ret));
+					return ERR_OCL_API;
+				}
 
-		if(xmrstak::params::inst().AMDCache)
+				std::ofstream file_stream;
+				file_stream.open(cache_file, std::ofstream::out | std::ofstream::binary);
+				file_stream.write(all_programs[dev_id], binary_sizes[dev_id]);
+				file_stream.close();
+				printer::inst()->print_msg(L1, "OpenCL device %u - Precompiled code stored in file %s",ctx->deviceIdx, cache_file.c_str());
+			}
+		}
+		else
 		{
-			int p_id = 0;
-			size_t mem_size = 0;
-			// create memory  structure to query all OpenCL program binaries
-			for(auto & p : all_programs)
+			printer::inst()->print_msg(L1, "OpenCL device %u - Load precompiled code from file %s",ctx->deviceIdx, cache_file.c_str());
+			std::ostringstream ss;
+			ss << clBinFile.rdbuf();
+			std::string s = ss.str();
+
+			size_t bin_size = s.size();
+			auto data_ptr = s.data();
+
+			cl_int clStatus;
+			ctx->Program[ii] = clCreateProgramWithBinary(
+				opencl_ctx, 1, &ctx->DeviceID, &bin_size,
+				(const unsigned char **)&data_ptr, &clStatus, &ret
+			);
+			if(ret != CL_SUCCESS)
 			{
-				program_storage.emplace_back(std::vector<char>(binary_sizes[p_id]));
-				all_programs[p_id] = program_storage[p_id].data();
-				mem_size += binary_sizes[p_id];
-				p_id++;
+				printer::inst()->print_msg(L1,"Error %s when calling clCreateProgramWithBinary. Try to delete file %s", err_to_str(ret), cache_file.c_str());
+				return ERR_OCL_API;
 			}
-
-		if((ret = clGetProgramInfo(ctx->Program, CL_PROGRAM_BINARIES, num_devices * sizeof(char*), all_programs.data(),NULL)) != CL_SUCCESS)
+			ret = clBuildProgram(ctx->Program[ii], 1, &ctx->DeviceID, NULL, NULL, NULL);
+			if(ret != CL_SUCCESS)
 			{
-				printer::inst()->print_msg(L1,"Error %s when calling clGetProgramInfo.", err_to_str(ret));
+				printer::inst()->print_msg(L1,"Error %s when calling clBuildProgram. Try to delete file %s", err_to_str(ret), cache_file.c_str());
 				return ERR_OCL_API;
 			}
-
-			std::ofstream file_stream;
-			file_stream.open(cache_file, std::ofstream::out | std::ofstream::binary);
-			file_stream.write(all_programs[dev_id], binary_sizes[dev_id]);
-			file_stream.close();
-			printer::inst()->print_msg(L1, "OpenCL device %u - Precompiled code stored in file %s",ctx->deviceIdx, cache_file.c_str());
 		}
-	}
-	else
-	{
-		printer::inst()->print_msg(L1, "OpenCL device %u - Load precompiled cod from file %s",ctx->deviceIdx, cache_file.c_str());
-		std::ostringstream ss;
-		ss << clBinFile.rdbuf();
-		std::string s = ss.str();
-
-		size_t bin_size = s.size();
-		auto data_ptr = s.data();
-
-		cl_int clStatus;
-		ctx->Program = clCreateProgramWithBinary(
-			opencl_ctx, 1, &ctx->DeviceID, &bin_size,
-			(const unsigned char **)&data_ptr, &clStatus, &ret
-		);
-		if(ret != CL_SUCCESS)
+
+		std::vector<std::string> KernelNames = { "cn0", "cn1", "cn2", "Blake", "Groestl", "JH", "Skein" };
+		// append algorithm number to kernel name
+		for(int k = 0; k < 3; k++)
+			KernelNames[k] += std::to_string(miner_algo[ii]);
+
+		if(ii == 0)
 		{
-			printer::inst()->print_msg(L1,"Error %s when calling clCreateProgramWithBinary. Try to delete file %s", err_to_str(ret), cache_file.c_str());
-			return ERR_OCL_API;
+			for(int i = 0; i < 7; ++i)
+			{
+				ctx->Kernels[ii][i] = clCreateKernel(ctx->Program[ii], KernelNames[i].c_str(), &ret);
+				if(ret != CL_SUCCESS)
+				{
+					printer::inst()->print_msg(L1,"Error %s when calling clCreateKernel for kernel_0 %s.", err_to_str(ret), KernelNames[i].c_str());
+					return ERR_OCL_API;
+				}
+			}
 		}
-		ret = clBuildProgram(ctx->Program, 1, &ctx->DeviceID, NULL, NULL, NULL);
-		if(ret != CL_SUCCESS)
+		else
 		{
-			printer::inst()->print_msg(L1,"Error %s when calling clBuildProgram. Try to delete file %s", err_to_str(ret), cache_file.c_str());
-			return ERR_OCL_API;
-		}
-	}
+			for(int i = 0; i < 3; ++i)
+			{
+				ctx->Kernels[ii][i] = clCreateKernel(ctx->Program[ii], KernelNames[i].c_str(), &ret);
+				if(ret != CL_SUCCESS)
+				{
+					printer::inst()->print_msg(L1,"Error %s when calling clCreateKernel for kernel_1 %s.", err_to_str(ret), KernelNames[i].c_str());
+					return ERR_OCL_API;
+				}
+			}
+			// move kernel from the main algorithm into the root algorithm kernel space
+			for(int i = 3; i < 7; ++i)
+			{
+				ctx->Kernels[ii][i] = ctx->Kernels[0][i];
+			}
 
-	const char *KernelNames[] = { "cn0", "cn1", "cn2", "Blake", "Groestl", "JH", "Skein", "cn1_monero" };
-	for(int i = 0; i < 8; ++i)
-	{
-		ctx->Kernels[i] = clCreateKernel(ctx->Program, KernelNames[i], &ret);
-		if(ret != CL_SUCCESS)
-		{
-			printer::inst()->print_msg(L1,"Error %s when calling clCreateKernel for kernel %s.", err_to_str(ret), KernelNames[i]);
-			return ERR_OCL_API;
 		}
 	}
-
 	ctx->Nonce = 0;
 	return 0;
 }
@@ -637,11 +675,18 @@ std::vector<GpuContext> getAMDDevices(int index)
 		}
 
 		std::string devVendor(devVendorVec.data());
-		if( devVendor.find("Advanced Micro Devices") != std::string::npos || devVendor.find("AMD") != std::string::npos)
+
+		bool isAMDDevice = devVendor.find("Advanced Micro Devices") != std::string::npos || devVendor.find("AMD") != std::string::npos;
+		bool isNVIDIADevice = devVendor.find("NVIDIA Corporation") != std::string::npos || devVendor.find("NVIDIA") != std::string::npos;
+
+		std::string selectedOpenCLVendor = xmrstak::params::inst().openCLVendor;
+		if((isAMDDevice && selectedOpenCLVendor == "AMD") || (isNVIDIADevice && selectedOpenCLVendor == "NVIDIA"))
 		{
 			GpuContext ctx;
 			std::vector<char> devNameVec(1024);
 			size_t maxMem;
+			if( devVendor.find("NVIDIA Corporation") != std::string::npos)
+				ctx.isNVIDIA = true;
 
 			if((clStatus = clGetDeviceInfo(device_list[k], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(int), &(ctx.computeUnits), NULL)) != CL_SUCCESS)
 			{
@@ -661,6 +706,10 @@ std::vector<GpuContext> getAMDDevices(int index)
 				continue;
 			}
 
+			// the allocation for NVIDIA OpenCL is not limited to 1/4 of the GPU memory per allocation
+			if(ctx.isNVIDIA)
+				maxMem = ctx.freeMem;
+
 			if((clStatus = clGetDeviceInfo(device_list[k], CL_DEVICE_NAME, devNameVec.size(), devNameVec.data(), NULL)) != CL_SUCCESS)
 			{
 				printer::inst()->print_msg(L1,"WARNING: %s when calling clGetDeviceInfo to get CL_DEVICE_NAME for device %u.", err_to_str(clStatus), k);
@@ -709,13 +758,15 @@ 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 ||
+
+			bool isAMDOpenCL = platformName.find("Advanced Micro Devices") != std::string::npos ||
 				platformName.find("Apple") != std::string::npos ||
-				platformName.find("Mesa") != std::string::npos
-			)
+				platformName.find("Mesa") != std::string::npos;
+			bool isNVIDIADevice = platformName.find("NVIDIA Corporation") != std::string::npos || platformName.find("NVIDIA") != std::string::npos;
+			std::string selectedOpenCLVendor = xmrstak::params::inst().openCLVendor;
+			if((isAMDOpenCL && selectedOpenCLVendor == "AMD") || (isNVIDIADevice && selectedOpenCLVendor == "NVIDIA"))
 			{
-
-				printer::inst()->print_msg(L0,"Found AMD platform index id = %i, name = %s",i , platformName.c_str());
+				printer::inst()->print_msg(L0,"Found %s platform index id = %i, name = %s", selectedOpenCLVendor.c_str(), i , platformName.c_str());
 				if(platformName.find("Mesa") != std::string::npos)
 					mesaPlatform = i;
 				else
@@ -781,7 +832,7 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx)
 	std::vector<char> platformNameVec(infoSize);
 	clGetPlatformInfo(PlatformIDList[platform_idx], CL_PLATFORM_VENDOR, infoSize, platformNameVec.data(), NULL);
 	std::string platformName(platformNameVec.data());
-	if( platformName.find("Advanced Micro Devices") == std::string::npos)
+	if(xmrstak::params::inst().openCLVendor == "AMD" && platformName.find("Advanced Micro Devices") == std::string::npos)
 	{
 		printer::inst()->print_msg(L1,"WARNING: using non AMD device: %s", platformName.c_str());
 	}
@@ -869,7 +920,8 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx)
 		{
 			size_t reduced_intensity = (ctx[i].rawIntensity / ctx[i].workSize) * ctx[i].workSize;
 			ctx[i].rawIntensity = reduced_intensity;
-			printer::inst()->print_msg(L0, "WARNING AMD: gpu %d intensity is not a multiple of 'worksize', auto reduce intensity to %d", ctx[i].deviceIdx, int(reduced_intensity));
+			const std::string backendName = xmrstak::params::inst().openCLVendor;
+			printer::inst()->print_msg(L0, "WARNING %s: gpu %d intensity is not a multiple of 'worksize', auto reduce intensity to %d", backendName.c_str(), ctx[i].deviceIdx, int(reduced_intensity));
 		}
 
 		if((ret = InitOpenCLGpu(opencl_ctx, &ctx[i], source_code.c_str())) != ERR_SUCCESS)
@@ -881,8 +933,11 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx)
 	return ERR_SUCCESS;
 }
 
-size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target, xmrstak_algo miner_algo, uint32_t version)
+size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target, xmrstak_algo miner_algo)
 {
+	// switch to the kernel storage 
+	int kernel_storage = miner_algo == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() ? 0 : 1;
+
 	cl_int ret;
 
 	if(input_len > 84)
@@ -899,71 +954,51 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
 		return ERR_OCL_API;
 	}
 
-	if((ret = clSetKernelArg(ctx->Kernels[0], 0, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS)
+	if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][0], 0, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 0, argument 0.", err_to_str(ret));
 		return ERR_OCL_API;
 	}
 
 	// Scratchpads
-	if((ret = clSetKernelArg(ctx->Kernels[0], 1, sizeof(cl_mem), ctx->ExtraBuffers + 0)) != CL_SUCCESS)
+	if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][0], 1, sizeof(cl_mem), ctx->ExtraBuffers + 0)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 0, argument 1.", err_to_str(ret));
 		return ERR_OCL_API;
 	}
 
 	// States
-	if((ret = clSetKernelArg(ctx->Kernels[0], 2, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS)
+	if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][0], 2, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 0, argument 2.", err_to_str(ret));
 		return ERR_OCL_API;
 	}
 
 	// Threads
-	if((ret = clSetKernelArg(ctx->Kernels[0], 3, sizeof(cl_ulong), &numThreads)) != CL_SUCCESS)
+	if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][0], 3, sizeof(cl_ulong), &numThreads)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 0, argument 3.", err_to_str(ret));
 		return(ERR_OCL_API);
 	}
 
-	/* ATTENTION: if we miner cryptonight_heavy the kernel needs an additional parameter version.
-	 * Do NOT use the variable `miner_algo` because this variable is changed dynamicly
-	 */
-	if(::jconf::inst()->GetMiningAlgo() == cryptonight_heavy)
-	{
-		// version
-		if ((ret = clSetKernelArg(ctx->Kernels[0], 4, sizeof(cl_uint), &version)) != CL_SUCCESS)
-		{
-			printer::inst()->print_msg(L1, "Error %s when calling clSetKernelArg for kernel 0, argument 4.", err_to_str(ret));
-			return ERR_OCL_API;
-		}
-	}
-
 	// CN1 Kernel
 
-	/// @todo only activate if currency is monero
-	int cn_kernel_offset = 0;
-	if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon)
-	{
-		cn_kernel_offset = 6;
-	}
-
 	// Scratchpads
-	if((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 0, sizeof(cl_mem), ctx->ExtraBuffers + 0)) != CL_SUCCESS)
+	if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][1], 0, sizeof(cl_mem), ctx->ExtraBuffers + 0)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 1, argument 0.", err_to_str(ret));
 		return ERR_OCL_API;
 	}
 
 	// States
-	if((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 1, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS)
+	if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][1], 1, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 1, argument 1.", err_to_str(ret));
 		return ERR_OCL_API;
 	}
 
 	// Threads
-	if((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 2, sizeof(cl_ulong), &numThreads)) != CL_SUCCESS)
+	if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][1], 2, sizeof(cl_ulong), &numThreads)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 1, argument 2.", err_to_str(ret));
 		return(ERR_OCL_API);
@@ -972,113 +1007,88 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
 	if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon )
 	{
 		// Input
-		if ((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS)
+		if ((ret = clSetKernelArg(ctx->Kernels[kernel_storage][1], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS)
 		{
 			printer::inst()->print_msg(L1, "Error %s when calling clSetKernelArg for kernel 1, arugment 4(input buffer).", err_to_str(ret));
 			return ERR_OCL_API;
 		}
 	}
-	/* ATTENTION: if we miner cryptonight_heavy the kernel needs an additional parameter version.
-	 * Do NOT use the variable `miner_algo` because this variable is changed dynamicly
-	 */
-	else if(::jconf::inst()->GetMiningAlgo() == cryptonight_heavy)
-	{
-		// version
-		if ((ret = clSetKernelArg(ctx->Kernels[1], 3, sizeof(cl_uint), &version)) != CL_SUCCESS)
-		{
-			printer::inst()->print_msg(L1, "Error %s when calling clSetKernelArg for kernel 1, argument 3 (version).", err_to_str(ret));
-			return ERR_OCL_API;
-		}
-	}
 
 	// CN3 Kernel
 	// Scratchpads
-	if((ret = clSetKernelArg(ctx->Kernels[2], 0, sizeof(cl_mem), ctx->ExtraBuffers + 0)) != CL_SUCCESS)
+	if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][2], 0, sizeof(cl_mem), ctx->ExtraBuffers + 0)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 2, argument 0.", err_to_str(ret));
 		return ERR_OCL_API;
 	}
 
 	// States
-	if((ret = clSetKernelArg(ctx->Kernels[2], 1, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS)
+	if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][2], 1, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 2, argument 1.", err_to_str(ret));
 		return ERR_OCL_API;
 	}
 
 	// Branch 0
-	if((ret = clSetKernelArg(ctx->Kernels[2], 2, sizeof(cl_mem), ctx->ExtraBuffers + 2)) != CL_SUCCESS)
+	if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][2], 2, sizeof(cl_mem), ctx->ExtraBuffers + 2)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 2, argument 2.", err_to_str(ret));
 		return ERR_OCL_API;
 	}
 
 	// Branch 1
-	if((ret = clSetKernelArg(ctx->Kernels[2], 3, sizeof(cl_mem), ctx->ExtraBuffers + 3)) != CL_SUCCESS)
+	if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][2], 3, sizeof(cl_mem), ctx->ExtraBuffers + 3)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 2, argument 3.", err_to_str(ret));
 		return ERR_OCL_API;
 	}
 
 	// Branch 2
-	if((ret = clSetKernelArg(ctx->Kernels[2], 4, sizeof(cl_mem), ctx->ExtraBuffers + 4)) != CL_SUCCESS)
+	if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][2], 4, sizeof(cl_mem), ctx->ExtraBuffers + 4)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 2, argument 4.", err_to_str(ret));
 		return ERR_OCL_API;
 	}
 
 	// Branch 3
-	if((ret = clSetKernelArg(ctx->Kernels[2], 5, sizeof(cl_mem), ctx->ExtraBuffers + 5)) != CL_SUCCESS)
+	if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][2], 5, sizeof(cl_mem), ctx->ExtraBuffers + 5)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 2, argument 5.", err_to_str(ret));
 		return ERR_OCL_API;
 	}
 
 	// Threads
-	if((ret = clSetKernelArg(ctx->Kernels[2], 6, sizeof(cl_ulong), &numThreads)) != CL_SUCCESS)
+	if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][2], 6, sizeof(cl_ulong), &numThreads)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel 2, argument 6.", err_to_str(ret));
 		return(ERR_OCL_API);
 	}
 
-    /* ATTENTION: if we miner cryptonight_heavy the kernel needs an additional parameter version.
-	 * Do NOT use the variable `miner_algo` because this variable is changed dynamicly
-	 */
-	if(::jconf::inst()->GetMiningAlgo() == cryptonight_heavy)
-	{
-		// version
-		if ((ret = clSetKernelArg(ctx->Kernels[2], 7, sizeof(cl_uint), &version)) != CL_SUCCESS)
-		{
-			printer::inst()->print_msg(L1, "Error %s when calling clSetKernelArg for kernel 2, argument 7.", err_to_str(ret));
-			return ERR_OCL_API;
-		}
-	}
-
 	for(int i = 0; i < 4; ++i)
 	{
 		// States
-		if((ret = clSetKernelArg(ctx->Kernels[i + 3], 0, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS)
+		if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][i + 3], 0, sizeof(cl_mem), ctx->ExtraBuffers + 1)) != CL_SUCCESS)
 		{
 			printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel %d, argument %d.", err_to_str(ret), i + 3, 0);
 			return ERR_OCL_API;
 		}
 
 		// Nonce buffer
-		if((ret = clSetKernelArg(ctx->Kernels[i + 3], 1, sizeof(cl_mem), ctx->ExtraBuffers + (i + 2))) != CL_SUCCESS)
+		if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][i + 3], 1, sizeof(cl_mem), ctx->ExtraBuffers + (i + 2))) != CL_SUCCESS)
 		{
 			printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel %d, argument %d.", err_to_str(ret), i + 3, 1);
 			return ERR_OCL_API;
 		}
 
 		// Output
-		if((ret = clSetKernelArg(ctx->Kernels[i + 3], 2, sizeof(cl_mem), &ctx->OutputBuffer)) != CL_SUCCESS)
+		if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][i + 3], 2, sizeof(cl_mem), &ctx->OutputBuffer)) != CL_SUCCESS)
 		{
 			printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel %d, argument %d.", err_to_str(ret), i + 3, 2);
 			return ERR_OCL_API;
 		}
 
 		// Target
-		if((ret = clSetKernelArg(ctx->Kernels[i + 3], 3, sizeof(cl_ulong), &target)) != CL_SUCCESS)
+		if((ret = clSetKernelArg(ctx->Kernels[kernel_storage][i + 3], 3, sizeof(cl_ulong), &target)) != CL_SUCCESS)
 		{
 			printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel %d, argument %d.", err_to_str(ret), i + 3, 3);
 			return ERR_OCL_API;
@@ -1088,8 +1098,11 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
 	return ERR_SUCCESS;
 }
 
-size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo, uint32_t version)
+size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo)
 {
+	// switch to the kernel storage
+	int kernel_storage = miner_algo == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() ? 0 : 1;
+	
 	cl_int ret;
 	cl_uint zero = 0;
 	size_t BranchNonces[4];
@@ -1125,35 +1138,21 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo,
 	clFinish(ctx->CommandQueues);
 
 	size_t Nonce[2] = {ctx->Nonce, 1}, gthreads[2] = { g_thd, 8 }, lthreads[2] = { w_size, 8 };
-	if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[0], 2, Nonce, gthreads, lthreads, 0, NULL, NULL)) != CL_SUCCESS)
+	if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[kernel_storage][0], 2, Nonce, gthreads, lthreads, 0, NULL, NULL)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clEnqueueNDRangeKernel for kernel %d.", err_to_str(ret), 0);
 		return ERR_OCL_API;
 	}
 
-	/*for(int i = 1; i < 3; ++i)
-	{
-		if((ret = clEnqueueNDRangeKernel(*ctx->CommandQueues, ctx->Kernels[i], 1, &ctx->Nonce, &g_thd, &w_size, 0, NULL, NULL)) != CL_SUCCESS)
-		{
-			Log(LOG_CRITICAL, "Error %s when calling clEnqueueNDRangeKernel for kernel %d.", err_to_str(ret), i);
-			return(ERR_OCL_API);
-		}
-	}*/
-
 	size_t tmpNonce = ctx->Nonce;
-	/// @todo only activate if currency is monero
-	int cn_kernel_offset = 0;
-	if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon)
-	{
-		cn_kernel_offset = 6;
-	}
-	if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[1 + cn_kernel_offset], 1, &tmpNonce, &g_thd, &w_size, 0, NULL, NULL)) != CL_SUCCESS)
+
+	if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[kernel_storage][1], 1, &tmpNonce, &g_thd, &w_size, 0, NULL, NULL)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clEnqueueNDRangeKernel for kernel %d.", err_to_str(ret), 1);
 		return ERR_OCL_API;
 	}
 
-	if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[2], 2, Nonce, gthreads, lthreads, 0, NULL, NULL)) != CL_SUCCESS)
+	if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[kernel_storage][2], 2, Nonce, gthreads, lthreads, 0, NULL, NULL)) != CL_SUCCESS)
 	{
 		printer::inst()->print_msg(L1,"Error %s when calling clEnqueueNDRangeKernel for kernel %d.", err_to_str(ret), 2);
 		return ERR_OCL_API;
@@ -1190,7 +1189,7 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo,
 		if(BranchNonces[i])
 		{
 			// Threads
-			if((clSetKernelArg(ctx->Kernels[i + 3], 4, sizeof(cl_ulong), BranchNonces + i)) != CL_SUCCESS)
+			if((clSetKernelArg(ctx->Kernels[kernel_storage][i + 3], 4, sizeof(cl_ulong), BranchNonces + i)) != CL_SUCCESS)
 			{
 				printer::inst()->print_msg(L1,"Error %s when calling clSetKernelArg for kernel %d, argument %d.", err_to_str(ret), i + 3, 4);
 				return(ERR_OCL_API);
@@ -1201,7 +1200,7 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo,
 			// number of global threads must be a multiple of the work group size (w_size)
 			assert(BranchNonces[i]%w_size == 0);
 			size_t tmpNonce = ctx->Nonce;
-			if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[i + 3], 1, &tmpNonce, BranchNonces + i, &w_size, 0, NULL, NULL)) != CL_SUCCESS)
+			if((ret = clEnqueueNDRangeKernel(ctx->CommandQueues, ctx->Kernels[kernel_storage][i + 3], 1, &tmpNonce, BranchNonces + i, &w_size, 0, NULL, NULL)) != CL_SUCCESS)
 			{
 				printer::inst()->print_msg(L1,"Error %s when calling clEnqueueNDRangeKernel for kernel %d.", err_to_str(ret), i + 3);
 				return ERR_OCL_API;
diff --git a/xmrstak/backend/amd/amd_gpu/gpu.hpp b/xmrstak/backend/amd/amd_gpu/gpu.hpp
index a387b15..5ab80b8 100644
--- a/xmrstak/backend/amd/amd_gpu/gpu.hpp
+++ b/xmrstak/backend/amd/amd_gpu/gpu.hpp
@@ -27,6 +27,7 @@ struct GpuContext
 	size_t workSize;
 	int stridedIndex;
 	int memChunk;
+	bool isNVIDIA = false;
 	int compMode;
 
 	/*Output vars*/
@@ -35,8 +36,8 @@ struct GpuContext
 	cl_mem InputBuffer;
 	cl_mem OutputBuffer;
 	cl_mem ExtraBuffers[6];
-	cl_program Program;
-	cl_kernel Kernels[8];
+	cl_program Program[2];
+	cl_kernel Kernels[2][8];
 	size_t freeMem;
 	int computeUnits;
 	std::string name;
@@ -50,7 +51,7 @@ int getAMDPlatformIdx();
 std::vector<GpuContext> getAMDDevices(int index);
 
 size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx);
-size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target, xmrstak_algo miner_algo, uint32_t version);
-size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo, uint32_t version);
+size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target, xmrstak_algo miner_algo);
+size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo);
 
 
diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
index 5d4e66c..d2ae1a7 100644
--- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
+++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
@@ -433,15 +433,13 @@ inline ulong getIdx()
 #endif
 }
 
-#define  mix_and_propagate(xin) (xin)[(get_local_id(1)) % 8][get_local_id(0)] ^ (xin)[(get_local_id(1) + 1) % 8][get_local_id(0)]
+#define mix_and_propagate(xin) (xin)[(get_local_id(1)) % 8][get_local_id(0)] ^ (xin)[(get_local_id(1) + 1) % 8][get_local_id(0)]
+		
+#define JOIN_DO(x,y) x##y
+#define JOIN(x,y) JOIN_DO(x,y)
 
 __attribute__((reqd_work_group_size(WORKSIZE, 8, 1)))
-__kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states, ulong Threads
-// cryptonight_heavy
-#if (ALGO == 4)
-		, uint version
-#endif
-)
+__kernel void JOIN(cn0,ALGO)(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states, ulong Threads)
 {
 	ulong State[25];
 	uint ExpandedKey1[40];
@@ -517,23 +515,20 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul
 		
 // cryptonight_heavy
 #if (ALGO == 4)
-	if(version >= 3)
-	{
-		__local uint4 xin[8][WORKSIZE];
+	__local uint4 xin[8][WORKSIZE];
 
-		/* Also left over threads performe this loop.
-		 * The left over thread results will be ignored
-		 */
-		for(size_t i=0; i < 16; i++)
-		{
-			#pragma unroll
-			for(int j = 0; j < 10; ++j)
-				text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey1)[j]);
-			barrier(CLK_LOCAL_MEM_FENCE);
-			xin[get_local_id(1)][get_local_id(0)] = text;
-			barrier(CLK_LOCAL_MEM_FENCE);
-			text = mix_and_propagate(xin);
-		}
+	/* Also left over threads performe this loop.
+	 * The left over thread results will be ignored
+	 */
+	for(size_t i=0; i < 16; i++)
+	{
+		#pragma unroll
+		for(int j = 0; j < 10; ++j)
+			text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey1)[j]);
+		barrier(CLK_LOCAL_MEM_FENCE);
+		xin[get_local_id(1)][get_local_id(0)] = text;
+		barrier(CLK_LOCAL_MEM_FENCE);
+		text = mix_and_propagate(xin);
 	}
 #endif
 
@@ -542,13 +537,9 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul
 	if(gIdx < Threads)
 #endif
 	{
-		int iterations = MEMORY >> 7;
-#if (ALGO == 4)
-		if(version < 3)
-			iterations >>= 1;
-#endif
+
 		#pragma unroll 2
-		for(int i = 0; i < iterations; ++i)
+		for(int i = 0; i < (MEMORY >> 7); ++i)
 		{
 			#pragma unroll
 			for(int j = 0; j < 10; ++j)
@@ -560,22 +551,13 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul
 	mem_fence(CLK_GLOBAL_MEM_FENCE);
 }
 
-#define VARIANT1_1(p) \
-	uint table = 0x75310U; \
-	uint index = (((p).s2 >> 26) & 12) | (((p).s2 >> 23) & 2); \
-	(p).s2 ^= ((table >> index) & 0x30U) << 24
-
-#define VARIANT1_2(p) ((uint2 *)&(p))[0] ^= tweak1_2
-
-#define VARIANT1_INIT() \
-	tweak1_2 = as_uint2(input[4]); \
-	tweak1_2.s0 >>= 24; \
-	tweak1_2.s0 |= tweak1_2.s1 << 8; \
-	tweak1_2.s1 = get_global_id(0); \
-	tweak1_2 ^= as_uint2(states[24])
-
 __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
-__kernel void cn1_monero(__global uint4 *Scratchpad, __global ulong *states, ulong Threads, __global ulong *input)
+__kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states, ulong Threads
+// cryptonight_monero || cryptonight_aeon
+#if(ALGO == 3 || ALGO == 5)
+, __global ulong *input
+#endif
+)
 {
 	ulong a[2], b[2];
 	__local uint AES0[256], AES1[256], AES2[256], AES3[256];
@@ -592,8 +574,9 @@ __kernel void cn1_monero(__global uint4 *Scratchpad, __global ulong *states, ulo
 	}
 
 	barrier(CLK_LOCAL_MEM_FENCE);
-
+#if(ALGO == 3 || ALGO == 5)
     uint2 tweak1_2;
+#endif
 	uint4 b_x;
 #if(COMP_MODE==1)
 	// do not use early return here
@@ -615,7 +598,13 @@ __kernel void cn1_monero(__global uint4 *Scratchpad, __global ulong *states, ulo
 		b[1] = states[3] ^ states[7];
 
 		b_x = ((uint4 *)b)[0];
-	    VARIANT1_INIT();
+#if(ALGO == 3 || ALGO == 5)
+		tweak1_2 = as_uint2(input[4]);
+		tweak1_2.s0 >>= 24;
+		tweak1_2.s0 |= tweak1_2.s1 << 8;
+		tweak1_2.s1 = get_global_id(0);
+		tweak1_2 ^= as_uint2(states[24]);
+#endif
 	}
 
 	mem_fence(CLK_LOCAL_MEM_FENCE);
@@ -625,17 +614,23 @@ __kernel void cn1_monero(__global uint4 *Scratchpad, __global ulong *states, ulo
 	if(gIdx < Threads)
 #endif
 	{
+		ulong idx0 = a[0];
+
 		#pragma unroll 8
 		for(int i = 0; i < ITERATIONS; ++i)
 		{
 			ulong c[2];
 
-			((uint4 *)c)[0] = Scratchpad[IDX((a[0] & MASK) >> 4)];
+			((uint4 *)c)[0] = Scratchpad[IDX((idx0 & MASK) >> 4)];
 			((uint4 *)c)[0] = AES_Round(AES0, AES1, AES2, AES3, ((uint4 *)c)[0], ((uint4 *)a)[0]);
 
 			b_x ^= ((uint4 *)c)[0];
-			VARIANT1_1(b_x);
-			Scratchpad[IDX((a[0] & MASK) >> 4)] = b_x;
+#if(ALGO == 3 || ALGO == 5)
+			uint table = 0x75310U;
+			uint index = ((b_x.s2 >> 26) & 12) | ((b_x.s2 >> 23) & 2);
+			b_x.s2 ^= ((table >> index) & 0x30U) << 24;
+#endif
+			Scratchpad[IDX((idx0 & MASK) >> 4)] = b_x;
 
 			uint4 tmp;
 			tmp = Scratchpad[IDX((c[0] & MASK) >> 4)];
@@ -643,101 +638,14 @@ __kernel void cn1_monero(__global uint4 *Scratchpad, __global ulong *states, ulo
 			a[1] += c[0] * as_ulong2(tmp).s0;
 			a[0] += mul_hi(c[0], as_ulong2(tmp).s0);
 
-			VARIANT1_2(a[1]);
-			Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0];
-			VARIANT1_2(a[1]);
-
-			((uint4 *)a)[0] ^= tmp;
-
-			b_x = ((uint4 *)c)[0];
-		}
-	}
-	mem_fence(CLK_GLOBAL_MEM_FENCE);
-}
-
-__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
-__kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Threads
-// cryptonight_heavy
-#if (ALGO == 4)
-		, uint version
-#endif
-)
-{
-	ulong a[2], b[2];
-	__local uint AES0[256], AES1[256], AES2[256], AES3[256];
-
-	const ulong gIdx = getIdx();
-
-	for(int i = get_local_id(0); i < 256; i += WORKSIZE)
-	{
-		const uint tmp = AES0_C[i];
-		AES0[i] = tmp;
-		AES1[i] = rotate(tmp, 8U);
-		AES2[i] = rotate(tmp, 16U);
-		AES3[i] = rotate(tmp, 24U);
-	}
-
-	barrier(CLK_LOCAL_MEM_FENCE);
-
-	uint4 b_x;
-#if(COMP_MODE==1)
-	// do not use early return here
-	if(gIdx < Threads)
-#endif
-	{
-		states += 25 * gIdx;
-#if(STRIDED_INDEX==0)
-		Scratchpad += gIdx * (MEMORY >> 4);
-#elif(STRIDED_INDEX==1)
-		Scratchpad += gIdx;
-#elif(STRIDED_INDEX==2)
-		Scratchpad += get_group_id(0) * (MEMORY >> 4) * WORKSIZE + MEM_CHUNK * get_local_id(0);
-#endif
-
-		a[0] = states[0] ^ states[4];
-		b[0] = states[2] ^ states[6];
-		a[1] = states[1] ^ states[5];
-		b[1] = states[3] ^ states[7];
-
-		b_x = ((uint4 *)b)[0];
-	}
-
-	mem_fence(CLK_LOCAL_MEM_FENCE);
-
-#if(COMP_MODE==1)
-	// do not use early return here
-	if(gIdx < Threads)
-#endif
-	{
-		ulong idx0 = a[0];
-		ulong mask = MASK;
 
-		int iterations = ITERATIONS;
-#if (ALGO == 4)
-		if(version < 3)
-		{
-			iterations <<= 1;
-			mask -= 0x200000;
-		}
+#if(ALGO == 3 || ALGO == 5)
+			((uint2 *)&(a[1]))[0] ^= tweak1_2;
+			Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0];
+			((uint2 *)&(a[1]))[0] ^= tweak1_2;
+#else
+			Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0];
 #endif
-		#pragma unroll 8
-		for(int i = 0; i < iterations; ++i)
-		{
-			ulong c[2];
-
-			((uint4 *)c)[0] = Scratchpad[IDX((idx0 & mask) >> 4)];
-			((uint4 *)c)[0] = AES_Round(AES0, AES1, AES2, AES3, ((uint4 *)c)[0], ((uint4 *)a)[0]);
-			//b_x ^= ((uint4 *)c)[0];
-
-			Scratchpad[IDX((idx0 & mask) >> 4)] = b_x ^ ((uint4 *)c)[0];
-
-			uint4 tmp;
-			tmp = Scratchpad[IDX((c[0] & mask) >> 4)];
-
-			a[1] += c[0] * as_ulong2(tmp).s0;
-			a[0] += mul_hi(c[0], as_ulong2(tmp).s0);
-
-			Scratchpad[IDX((c[0] & mask) >> 4)] = ((uint4 *)a)[0];
 
 			((uint4 *)a)[0] ^= tmp;
 			idx0 = a[0];
@@ -745,14 +653,11 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre
 			b_x = ((uint4 *)c)[0];
 // cryptonight_heavy
 #if (ALGO == 4)
-			if(version >= 3)
-			{
-				long n = *((__global long*)(Scratchpad + (IDX((idx0 & mask) >> 4))));
-				int d = ((__global int*)(Scratchpad + (IDX((idx0 & mask) >> 4))))[2];
-				long q = n / (d | 0x5);
-				*((__global long*)(Scratchpad + (IDX((idx0 & mask) >> 4)))) = n ^ q;
-				idx0 = d ^ q;
-			}
+			long n = *((__global long*)(Scratchpad + (IDX((idx0 & MASK) >> 4))));
+			int d = ((__global int*)(Scratchpad + (IDX((idx0 & MASK) >> 4))))[2];
+			long q = n / (d | 0x5);
+			*((__global long*)(Scratchpad + (IDX((idx0 & MASK) >> 4)))) = n ^ q;
+			idx0 = d ^ q;
 #endif
 		}
 	}
@@ -760,12 +665,7 @@ __kernel void cn1(__global uint4 *Scratchpad, __global ulong *states, ulong Thre
 }
 
 __attribute__((reqd_work_group_size(WORKSIZE, 8, 1)))
-__kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global uint *Branch0, __global uint *Branch1, __global uint *Branch2, __global uint *Branch3, ulong Threads
-// cryptonight_heavy
-#if (ALGO == 4)
-	, uint version
-#endif
-		)
+__kernel void JOIN(cn2,ALGO) (__global uint4 *Scratchpad, __global ulong *states, __global uint *Branch0, __global uint *Branch1, __global uint *Branch2, __global uint *Branch3, ulong Threads)
 {
 	__local uint AES0[256], AES1[256], AES2[256], AES3[256];
 	uint ExpandedKey2[40];
@@ -827,58 +727,42 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u
 	if(gIdx < Threads)
 #endif
 	{
-		int iterations = MEMORY >> 7;
 #if (ALGO == 4)
-		if(version < 3)
-		{
-			iterations >>= 1;
-			#pragma unroll 2
-			for(int i = 0; i < iterations; ++i)
-			{
-				text ^= Scratchpad[IDX((i << 3) + get_local_id(1))];
-
-				#pragma unroll 10
-				for(int j = 0; j < 10; ++j)
-					text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]);
-			}
-		}
-		else
+		#pragma unroll 2
+		for(int i = 0; i < (MEMORY >> 7); ++i)
 		{
-			#pragma unroll 2
-			for(int i = 0; i < iterations; ++i)
-			{
-				text ^= Scratchpad[IDX((i << 3) + get_local_id(1))];
+			text ^= Scratchpad[IDX((i << 3) + get_local_id(1))];
 
-				#pragma unroll 10
-				for(int j = 0; j < 10; ++j)
-					text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]);
+			#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);
-				xin[get_local_id(1)][get_local_id(0)] = text;
-				barrier(CLK_LOCAL_MEM_FENCE);
-				text = mix_and_propagate(xin);
-			}
+			barrier(CLK_LOCAL_MEM_FENCE);
+			xin[get_local_id(1)][get_local_id(0)] = text;
+			barrier(CLK_LOCAL_MEM_FENCE);
+			text = mix_and_propagate(xin);
+		}
 
-			#pragma unroll 2
-			for(int i = 0; i < iterations; ++i)
-			{
-				text ^= Scratchpad[IDX((i << 3) + get_local_id(1))];
+		#pragma unroll 2
+		for(int i = 0; i < (MEMORY >> 7); ++i)
+		{
+			text ^= Scratchpad[IDX((i << 3) + get_local_id(1))];
 
-				#pragma unroll 10
-				for(int j = 0; j < 10; ++j)
-					text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]);
+			#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);
-				xin[get_local_id(1)][get_local_id(0)] = text;
-				barrier(CLK_LOCAL_MEM_FENCE);
-				text = mix_and_propagate(xin);
-			}
+			barrier(CLK_LOCAL_MEM_FENCE);
+			xin[get_local_id(1)][get_local_id(0)] = text;
+			barrier(CLK_LOCAL_MEM_FENCE);
+			text = mix_and_propagate(xin);
 		}
+		
 #else
 		#pragma unroll 2
-		for(int i = 0; i < iterations; ++i)
+		for(int i = 0; i < (MEMORY >> 7); ++i)
 		{
 			text ^= Scratchpad[IDX((i << 3) + get_local_id(1))];
 
@@ -891,21 +775,18 @@ __kernel void cn2(__global uint4 *Scratchpad, __global ulong *states, __global u
 
 // cryptonight_heavy
 #if (ALGO == 4)
-	if(version >= 3)
+	/* Also left over threads perform this loop.
+	 * The left over thread results will be ignored
+	 */
+	for(size_t i=0; i < 16; i++)
 	{
-		/* Also left over threads performe this loop.
-		 * The left over thread results will be ignored
-		 */
-		for(size_t i=0; i < 16; i++)
-		{
-			#pragma unroll
-			for(int j = 0; j < 10; ++j)
-				text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]);
-			barrier(CLK_LOCAL_MEM_FENCE);
-			xin[get_local_id(1)][get_local_id(0)] = text;
-			barrier(CLK_LOCAL_MEM_FENCE);
-			text = mix_and_propagate(xin);
-		}
+		#pragma unroll
+		for(int j = 0; j < 10; ++j)
+			text = AES_Round(AES0, AES1, AES2, AES3, text, ((uint4 *)ExpandedKey2)[j]);
+		barrier(CLK_LOCAL_MEM_FENCE);
+		xin[get_local_id(1)][get_local_id(0)] = text;
+		barrier(CLK_LOCAL_MEM_FENCE);
+		text = mix_and_propagate(xin);
 	}
 #endif
 
diff --git a/xmrstak/backend/amd/autoAdjust.hpp b/xmrstak/backend/amd/autoAdjust.hpp
index ea057a0..685890b 100644
--- a/xmrstak/backend/amd/autoAdjust.hpp
+++ b/xmrstak/backend/amd/autoAdjust.hpp
@@ -83,11 +83,15 @@ private:
 
 		constexpr size_t byteToMiB = 1024u * 1024u;
 
-		size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
+		size_t hashMemSize = std::max(
+			cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()),
+			cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot())
+		);
 
 		std::string conf;
 		for(auto& ctx : devVec)
 		{
+			size_t minFreeMem = 128u * byteToMiB;
 			/* 1000 is a magic selected limit, the reason is that more than 2GiB memory
 			 * sowing down the memory performance because of TLB cache misses
 			 */
@@ -109,12 +113,26 @@ private:
 				 */
 				maxThreads = 2024u;
 			}
+
+			// NVIDIA optimizations
+			if(
+				ctx.isNVIDIA && (
+					ctx.name.find("P100") != std::string::npos ||
+				    ctx.name.find("V100") != std::string::npos
+				)
+			)
+			{
+				// do not limit the number of threads
+				maxThreads = 40000u;
+				minFreeMem = 512u * byteToMiB;
+			}
+
 			// increase all intensity limits by two for aeon
-			if(::jconf::inst()->GetMiningAlgo() == cryptonight_lite)
+			if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_lite)
 				maxThreads *= 2u;
 
 			// keep 128MiB memory free (value is randomly chosen)
-			size_t availableMem = ctx.freeMem - (128u * byteToMiB);
+			size_t availableMem = ctx.freeMem - minFreeMem;
 			// 224byte extra memory is used per thread for meta data
 			size_t perThread = hashMemSize + 224u;
 			size_t maxIntensity = availableMem / perThread;
@@ -135,7 +153,7 @@ private:
 				// set 8 threads per block (this is a good value for the most gpus)
 				conf += std::string("  { \"index\" : ") + std::to_string(ctx.deviceIdx) + ",\n" +
 					"    \"intensity\" : " + std::to_string(intensity) + ", \"worksize\" : " + std::to_string(8) + ",\n" +
-					"    \"affine_to_cpu\" : false, \"strided_index\" : 1, \"mem_chunk\" : 2,\n"
+					"    \"affine_to_cpu\" : false, \"strided_index\" : " + (ctx.isNVIDIA ? "0" : "1") + ", \"mem_chunk\" : 2,\n"
 					"    \"comp_mode\" : true\n" +
 					"  },\n";
 			}
@@ -148,11 +166,13 @@ private:
 		configTpl.replace("PLATFORMINDEX",std::to_string(platformIndex));
 		configTpl.replace("GPUCONFIG",conf);
 		configTpl.write(params::inst().configFileAMD);
-		printer::inst()->print_msg(L0, "AMD: GPU configuration stored in file '%s'", params::inst().configFileAMD.c_str());
+
+		const std::string backendName = xmrstak::params::inst().openCLVendor;
+		printer::inst()->print_msg(L0, "%s: GPU (OpenCL) configuration stored in file '%s'", backendName.c_str(), params::inst().configFileAMD.c_str());
 	}
 
 	std::vector<GpuContext> devVec;
 };
 
 } // namespace amd
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp
index cab5ad9..4353e3d 100644
--- a/xmrstak/backend/amd/minethd.cpp
+++ b/xmrstak/backend/amd/minethd.cpp
@@ -137,6 +137,8 @@ std::vector<iBackend*>* minethd::thread_starter(uint32_t threadOffset, miner_wor
 	for (i = 0; i < n; i++)
 	{
 		jconf::inst()->GetThreadConfig(i, cfg);
+
+		const std::string backendName = xmrstak::params::inst().openCLVendor;
 		
 		if(cfg.cpu_aff >= 0)
 		{
@@ -144,10 +146,10 @@ std::vector<iBackend*>* minethd::thread_starter(uint32_t threadOffset, miner_wor
 			printer::inst()->print_msg(L1, "WARNING on macOS thread affinity is only advisory.");
 #endif
 
-			printer::inst()->print_msg(L1, "Starting AMD GPU thread %d, affinity: %d.", i, (int)cfg.cpu_aff);
+			printer::inst()->print_msg(L1, "Starting %s GPU (OpenCL) thread %d, affinity: %d.", backendName.c_str(), i, (int)cfg.cpu_aff);
 		}
 		else
-			printer::inst()->print_msg(L1, "Starting AMD GPU thread %d, no affinity.", i);
+			printer::inst()->print_msg(L1, "Starting %s GPU (OpenCL) thread %d, no affinity.", backendName.c_str(), i);
 
 		minethd* thd = new minethd(pWork, i + threadOffset, &vGpuData[i], cfg);
 		pvThreads->push_back(thd);
@@ -193,12 +195,13 @@ void minethd::work_main()
 	cpu_ctx = cpu::minethd::minethd_alloc_ctx();
 	
 	// start with root algorithm and switch later if fork version is reached
-	auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+	auto miner_algo = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot();
 	cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
 
 	globalStates::inst().iConsumeCnt++;
 
 	uint8_t version = 0;
+	size_t lastPoolId = 0;
 
 	while (bQuit == 0)
 	{
@@ -217,13 +220,20 @@ void minethd::work_main()
 		}
 
 		uint8_t new_version = oWork.getVersion();
-		if(new_version != version)
+		if(new_version != version || oWork.iPoolId != lastPoolId)
 		{
-			if(new_version >= ::jconf::inst()->GetMiningForkVersion())
+			coinDescription coinDesc = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(oWork.iPoolId);
+			if(new_version >= coinDesc.GetMiningForkVersion())
+			{
+				miner_algo = coinDesc.GetMiningAlgo();
+				hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
+			}
+			else
 			{
-				miner_algo = ::jconf::inst()->GetMiningAlgo();
+				miner_algo = coinDesc.GetMiningAlgoRoot();
 				hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
 			}
+			lastPoolId = oWork.iPoolId;
 			version = new_version;
 		}
 
@@ -233,7 +243,7 @@ void minethd::work_main()
 		assert(sizeof(job_result::sJobID) == sizeof(pool_job::sJobID));
 		uint64_t target = oWork.iTarget;
 		
-		XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target, miner_algo, version);
+		XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target, miner_algo);
 
 		if(oWork.bNiceHash)
 			pGpuCtx->Nonce = *(uint32_t*)(oWork.bWorkBlob + 39);
@@ -249,7 +259,7 @@ void minethd::work_main()
 			cl_uint results[0x100];
 			memset(results,0,sizeof(cl_uint)*(0x100));
 
-			XMRRunJob(pGpuCtx, results, miner_algo, version);
+			XMRRunJob(pGpuCtx, results, miner_algo);
 
 			for(size_t i = 0; i < results[0xFF]; i++)
 			{
diff --git a/xmrstak/backend/backendConnector.cpp b/xmrstak/backend/backendConnector.cpp
index d735cb3..6f80a0f 100644
--- a/xmrstak/backend/backendConnector.cpp
+++ b/xmrstak/backend/backendConnector.cpp
@@ -77,11 +77,12 @@ std::vector<iBackend*>* BackendConnector::thread_starter(miner_work& pWork)
 #ifndef CONF_NO_OPENCL
 	if(params::inst().useAMD)
 	{
-		plugin amdplugin("AMD", "xmrstak_opencl_backend");
+		const std::string backendName = xmrstak::params::inst().openCLVendor;
+		plugin amdplugin(backendName, "xmrstak_opencl_backend");
 		std::vector<iBackend*>* amdThreads = amdplugin.startBackend(static_cast<uint32_t>(pvThreads->size()), pWork, environment::inst());
 		pvThreads->insert(std::end(*pvThreads), std::begin(*amdThreads), std::end(*amdThreads));
 		if(amdThreads->size() == 0)
-			printer::inst()->print_msg(L0, "WARNING: backend AMD disabled.");
+			printer::inst()->print_msg(L0, "WARNING: backend %s (OpenCL) disabled.", backendName.c_str());
 	}
 #endif
 
@@ -99,4 +100,4 @@ std::vector<iBackend*>* BackendConnector::thread_starter(miner_work& pWork)
 	return pvThreads;
 }
 
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/backendConnector.hpp b/xmrstak/backend/backendConnector.hpp
index da3dc77..66d873e 100644
--- a/xmrstak/backend/backendConnector.hpp
+++ b/xmrstak/backend/backendConnector.hpp
@@ -18,4 +18,4 @@ namespace xmrstak
 		static bool self_test();
 	};
 
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/cpu/autoAdjust.hpp b/xmrstak/backend/cpu/autoAdjust.hpp
index 969d478..518721a 100644
--- a/xmrstak/backend/cpu/autoAdjust.hpp
+++ b/xmrstak/backend/cpu/autoAdjust.hpp
@@ -36,7 +36,10 @@ public:
 	bool printConfig()
 	{
 
-		const size_t hashMemSizeKB = cn_select_memory(::jconf::inst()->GetMiningAlgo()) / 1024u;
+		const size_t hashMemSizeKB = std::max(
+			cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()),
+			cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot())
+		) / 1024u;
 		const size_t halfHashMemSizeKB = hashMemSizeKB / 2u;
 
 		configEditor configTpl{};
@@ -172,4 +175,4 @@ private:
 };
 
 } // namespace cpu
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/cpu/autoAdjustHwloc.hpp b/xmrstak/backend/cpu/autoAdjustHwloc.hpp
index 568abb5..b1f3914 100644
--- a/xmrstak/backend/cpu/autoAdjustHwloc.hpp
+++ b/xmrstak/backend/cpu/autoAdjustHwloc.hpp
@@ -28,7 +28,10 @@ public:
 
 	autoAdjust()
 	{
-		hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
+		hashMemSize = std::max(
+			cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()),
+			cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot())
+		);
 		halfHashMemSize = hashMemSize / 2u;
 	}
 
@@ -214,4 +217,4 @@ private:
 };
 
 } // namespace cpu
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/cpu/crypto/c_blake256.c b/xmrstak/backend/cpu/crypto/c_blake256.c
index ff623dd..e5fadfe 100644
--- a/xmrstak/backend/cpu/crypto/c_blake256.c
+++ b/xmrstak/backend/cpu/crypto/c_blake256.c
@@ -124,7 +124,7 @@ void blake224_init(state *S) {
 }
 
 // datalen = number of bits
-void blake256_update(state *S, const uint8_t *data, uint64_t datalen) {
+void blake256_update(state *S, const uint8_t *data, uint32_t datalen) {
 	int left = S->buflen >> 3;
 	int fill = 64 - left;
 
@@ -155,7 +155,7 @@ void blake256_update(state *S, const uint8_t *data, uint64_t datalen) {
 }
 
 // datalen = number of bits
-void blake224_update(state *S, const uint8_t *data, uint64_t datalen) {
+void blake224_update(state *S, const uint8_t *data, uint32_t datalen) {
 	blake256_update(S, data, datalen);
 }
 
@@ -206,7 +206,7 @@ void blake224_final(state *S, uint8_t *digest) {
 }
 
 // inlen = number of bytes
-void blake256_hash(uint8_t *out, const uint8_t *in, uint64_t inlen) {
+void blake256_hash(uint8_t *out, const uint8_t *in, uint32_t inlen) {
 	state S;
 	blake256_init(&S);
 	blake256_update(&S, in, inlen * 8);
@@ -214,7 +214,7 @@ void blake256_hash(uint8_t *out, const uint8_t *in, uint64_t inlen) {
 }
 
 // inlen = number of bytes
-void blake224_hash(uint8_t *out, const uint8_t *in, uint64_t inlen) {
+void blake224_hash(uint8_t *out, const uint8_t *in, uint32_t inlen) {
 	state S;
 	blake224_init(&S);
 	blake224_update(&S, in, inlen * 8);
@@ -282,13 +282,13 @@ void hmac_blake224_init(hmac_state *S, const uint8_t *_key, uint64_t keylen) {
 }
 
 // datalen = number of bits
-void hmac_blake256_update(hmac_state *S, const uint8_t *data, uint64_t datalen) {
+void hmac_blake256_update(hmac_state *S, const uint8_t *data, uint32_t datalen) {
   // update the inner state
   blake256_update(&S->inner, data, datalen);
 }
 
 // datalen = number of bits
-void hmac_blake224_update(hmac_state *S, const uint8_t *data, uint64_t datalen) {
+void hmac_blake224_update(hmac_state *S, const uint8_t *data, uint32_t datalen) {
   // update the inner state
   blake224_update(&S->inner, data, datalen);
 }
@@ -310,7 +310,7 @@ void hmac_blake224_final(hmac_state *S, uint8_t *digest) {
 }
 
 // keylen = number of bytes; inlen = number of bytes
-void hmac_blake256_hash(uint8_t *out, const uint8_t *key, uint64_t keylen, const uint8_t *in, uint64_t inlen) {
+void hmac_blake256_hash(uint8_t *out, const uint8_t *key, uint64_t keylen, const uint8_t *in, uint32_t inlen) {
 	hmac_state S;
 	hmac_blake256_init(&S, key, keylen);
 	hmac_blake256_update(&S, in, inlen * 8);
@@ -318,7 +318,7 @@ void hmac_blake256_hash(uint8_t *out, const uint8_t *key, uint64_t keylen, const
 }
 
 // keylen = number of bytes; inlen = number of bytes
-void hmac_blake224_hash(uint8_t *out, const uint8_t *key, uint64_t keylen, const uint8_t *in, uint64_t inlen) {
+void hmac_blake224_hash(uint8_t *out, const uint8_t *key, uint64_t keylen, const uint8_t *in, uint32_t inlen) {
 	hmac_state S;
 	hmac_blake224_init(&S, key, keylen);
 	hmac_blake224_update(&S, in, inlen * 8);
diff --git a/xmrstak/backend/cpu/crypto/c_blake256.h b/xmrstak/backend/cpu/crypto/c_blake256.h
index b9c2aad..06c7917 100644
--- a/xmrstak/backend/cpu/crypto/c_blake256.h
+++ b/xmrstak/backend/cpu/crypto/c_blake256.h
@@ -17,27 +17,27 @@ typedef struct {
 void blake256_init(state *);
 void blake224_init(state *);
 
-void blake256_update(state *, const uint8_t *, uint64_t);
-void blake224_update(state *, const uint8_t *, uint64_t);
+void blake256_update(state *, const uint8_t *, uint32_t);
+void blake224_update(state *, const uint8_t *, uint32_t);
 
 void blake256_final(state *, uint8_t *);
 void blake224_final(state *, uint8_t *);
 
-void blake256_hash(uint8_t *, const uint8_t *, uint64_t);
-void blake224_hash(uint8_t *, const uint8_t *, uint64_t);
+void blake256_hash(uint8_t *, const uint8_t *, uint32_t);
+void blake224_hash(uint8_t *, const uint8_t *, uint32_t);
 
 /* HMAC functions: */
 
 void hmac_blake256_init(hmac_state *, const uint8_t *, uint64_t);
 void hmac_blake224_init(hmac_state *, const uint8_t *, uint64_t);
 
-void hmac_blake256_update(hmac_state *, const uint8_t *, uint64_t);
-void hmac_blake224_update(hmac_state *, const uint8_t *, uint64_t);
+void hmac_blake256_update(hmac_state *, const uint8_t *, uint32_t);
+void hmac_blake224_update(hmac_state *, const uint8_t *, uint32_t);
 
 void hmac_blake256_final(hmac_state *, uint8_t *);
 void hmac_blake224_final(hmac_state *, uint8_t *);
 
-void hmac_blake256_hash(uint8_t *, const uint8_t *, uint64_t, const uint8_t *, uint64_t);
-void hmac_blake224_hash(uint8_t *, const uint8_t *, uint64_t, const uint8_t *, uint64_t);
+void hmac_blake256_hash(uint8_t *, const uint8_t *, uint64_t, const uint8_t *, uint32_t);
+void hmac_blake224_hash(uint8_t *, const uint8_t *, uint64_t, const uint8_t *, uint32_t);
 
 #endif /* _BLAKE256_H_ */
diff --git a/xmrstak/backend/cpu/crypto/c_skein.h b/xmrstak/backend/cpu/crypto/c_skein.h
index 6165a2a..86dbc08 100644
--- a/xmrstak/backend/cpu/crypto/c_skein.h
+++ b/xmrstak/backend/cpu/crypto/c_skein.h
@@ -37,7 +37,7 @@ typedef enum
 }
 SkeinHashReturn;
 
-typedef size_t   SkeinDataLength;                /* bit count  type */
+typedef uint32_t SkeinDataLength;                /* bit count  type */
 typedef u08b_t   SkeinBitSequence;               /* bit stream type */
 
 /* "all-in-one" call */
diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
index 5203de8..7562de1 100644
--- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
+++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h
@@ -43,7 +43,7 @@ extern "C"
 {
 	void keccak(const uint8_t *in, int inlen, uint8_t *md, int mdlen);
 	void keccakf(uint64_t st[25], int rounds);
-	extern void(*const extra_hashes[4])(const void *, size_t, char *);
+	extern void(*const extra_hashes[4])(const void *, uint32_t, char *);
 }
 
 // This will shift and xor tmp1 into itself as 4 32-bit vals such as
@@ -429,7 +429,7 @@ inline void cryptonight_monero_tweak(uint64_t* mem_out, __m128i tmp)
 	tmp = _mm_castps_si128(_mm_movehl_ps(_mm_castsi128_ps(tmp), _mm_castsi128_ps(tmp)));
 	uint64_t vh = _mm_cvtsi128_si64(tmp);
 
-	uint8_t x = vh >> 24;
+	uint8_t x = static_cast<uint8_t>(vh >> 24);
 	static const uint16_t table = 0x7531;
 	const uint8_t index = (((x >> 3) & 6) | (x & 1)) << 1;
 	vh ^= ((table >> index) & 0x3) << 28;
diff --git a/xmrstak/backend/cpu/crypto/cryptonight_common.cpp b/xmrstak/backend/cpu/crypto/cryptonight_common.cpp
index 17fa24b..ee3b663 100644
--- a/xmrstak/backend/cpu/crypto/cryptonight_common.cpp
+++ b/xmrstak/backend/cpu/crypto/cryptonight_common.cpp
@@ -35,6 +35,7 @@ extern "C"
 #include "xmrstak/jconf.hpp"
 #include <stdio.h>
 #include <stdlib.h>
+#include <algorithm>
 
 #ifdef __GNUC__
 #include <mm_malloc.h>
@@ -55,23 +56,23 @@ extern "C"
 #include <string.h>
 #endif // _WIN32
 
-void do_blake_hash(const void* input, size_t len, char* output) {
+void do_blake_hash(const void* input, uint32_t len, char* output) {
 	blake256_hash((uint8_t*)output, (const uint8_t*)input, len);
 }
 
-void do_groestl_hash(const void* input, size_t len, char* output) {
+void do_groestl_hash(const void* input, uint32_t len, char* output) {
 	groestl((const uint8_t*)input, len * 8, (uint8_t*)output);
 }
 
-void do_jh_hash(const void* input, size_t len, char* output) {
+void do_jh_hash(const void* input, uint32_t len, char* output) {
 	jh_hash(32 * 8, (const uint8_t*)input, 8 * len, (uint8_t*)output);
 }
 
-void do_skein_hash(const void* input, size_t len, char* output) {
+void do_skein_hash(const void* input, uint32_t len, char* output) {
 	skein_hash(8 * 32, (const uint8_t*)input, 8 * len, (uint8_t*)output);
 }
 
-void (* const extra_hashes[4])(const void *, size_t, char *) = {do_blake_hash, do_groestl_hash, do_jh_hash, do_skein_hash};
+void (* const extra_hashes[4])(const void *, uint32_t, char *) = {do_blake_hash, do_groestl_hash, do_jh_hash, do_skein_hash};
 
 #ifdef _WIN32
 #include "xmrstak/misc/uac.hpp"
@@ -202,7 +203,10 @@ size_t cryptonight_init(size_t use_fast_mem, size_t use_mlock, alloc_msg* msg)
 
 cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, alloc_msg* msg)
 {
-	size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
+	size_t hashMemSize = std::max(
+		cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()),
+		cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot())
+	);
 
 	cryptonight_ctx* ptr = (cryptonight_ctx*)_mm_malloc(sizeof(cryptonight_ctx), 4096);
 
@@ -278,7 +282,10 @@ cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, al
 
 void cryptonight_free_ctx(cryptonight_ctx* ctx)
 {
-	size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
+	size_t hashMemSize = std::max(
+		cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()),
+		cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot())
+	);
 
 	if(ctx->ctx_info[0] != 0)
 	{
diff --git a/xmrstak/backend/cpu/crypto/hash.h b/xmrstak/backend/cpu/crypto/hash.h
index c12d355..2af3309 100644
--- a/xmrstak/backend/cpu/crypto/hash.h
+++ b/xmrstak/backend/cpu/crypto/hash.h
@@ -1,5 +1,7 @@
 #pragma once
 
+#include <stdint.h>
+
 typedef unsigned char BitSequence;
-typedef unsigned long long DataLength;
+typedef uint32_t DataLength;
 typedef enum {SUCCESS = 0, FAIL = 1, BAD_HASHLEN = 2} HashReturn;
diff --git a/xmrstak/backend/cpu/crypto/skein_port.h b/xmrstak/backend/cpu/crypto/skein_port.h
index 9cbefcb..99641bc 100644
--- a/xmrstak/backend/cpu/crypto/skein_port.h
+++ b/xmrstak/backend/cpu/crypto/skein_port.h
@@ -49,7 +49,7 @@
 								multiple of size / 8)
 
 	ptr_cast(x,size)            casts a pointer to a pointer to a
-								varaiable of length 'size' bits
+								variable of length 'size' bits
 */
 
 #define ui_type(size)               uint##size##_t
diff --git a/xmrstak/backend/cpu/hwlocMemory.cpp b/xmrstak/backend/cpu/hwlocMemory.cpp
index 94d2b53..089570f 100644
--- a/xmrstak/backend/cpu/hwlocMemory.cpp
+++ b/xmrstak/backend/cpu/hwlocMemory.cpp
@@ -30,7 +30,7 @@ void bindMemoryToNUMANode( size_t puId )
 
 	depth = hwloc_get_type_depth(topology, HWLOC_OBJ_PU);
 
-	for( size_t i = 0;
+	for( uint32_t i = 0;
 		i < hwloc_get_nbobjs_by_depth(topology, depth);
 		i++ )
 	{
diff --git a/xmrstak/backend/cpu/jconf.cpp b/xmrstak/backend/cpu/jconf.cpp
index 399dd16..49da7ae 100644
--- a/xmrstak/backend/cpu/jconf.cpp
+++ b/xmrstak/backend/cpu/jconf.cpp
@@ -259,4 +259,4 @@ bool jconf::parse_config(const char* sFilename)
 }
 
 } // namespace cpu
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/cpu/jconf.hpp b/xmrstak/backend/cpu/jconf.hpp
index e98ed16..be85503 100644
--- a/xmrstak/backend/cpu/jconf.hpp
+++ b/xmrstak/backend/cpu/jconf.hpp
@@ -40,4 +40,4 @@ private:
 };
 
 } // namespace cpu
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp
index 5b56f85..f8f70f9 100644
--- a/xmrstak/backend/cpu/minethd.cpp
+++ b/xmrstak/backend/cpu/minethd.cpp
@@ -231,7 +231,7 @@ bool minethd::self_test()
 
 	bool bResult = true;
 
-	if(::jconf::inst()->GetMiningAlgo() == cryptonight)
+	if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight)
 	{
 		unsigned char out[32 * MAX_N];
 		cn_hash_fun hashf;
@@ -276,13 +276,13 @@ bool minethd::self_test()
 				"\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05"
 				"\xa0\x84\xf0\x1d\x14\x37\xa0\x9c\x69\x85\x40\x1b\x60\xd4\x35\x54\xae\x10\x58\x02\xc5\xf5\xd8\xa9\xb3\x25\x36\x49\xc0\xbe\x66\x05", 160) == 0;
 	}
-	else if(::jconf::inst()->GetMiningAlgo() == cryptonight_lite)
+	else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_lite)
 	{
 	}
-	else if(::jconf::inst()->GetMiningAlgo() == cryptonight_monero)
+	else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_monero)
 	{
 	}
-	else if(::jconf::inst()->GetMiningAlgo() == cryptonight_aeon)
+	else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_aeon)
 	{
 	}
 
@@ -424,7 +424,7 @@ void minethd::work_main()
 	job_result result;
 
 	// start with root algorithm and switch later if fork version is reached
-	auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+	auto miner_algo = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot();
 	cn_hash_fun hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
 	ctx = minethd_alloc_ctx();
 
@@ -434,6 +434,7 @@ void minethd::work_main()
 	result.iThreadId = iThreadNo;
 
 	uint8_t version = 0;
+	size_t lastPoolId = 0;
 
 	while (bQuit == 0)
 	{
@@ -461,13 +462,20 @@ void minethd::work_main()
 			result.iNonce = *piNonce;
 
 		uint8_t new_version = oWork.getVersion();
-		if(new_version != version)
+		if(new_version != version || oWork.iPoolId != lastPoolId)
 		{
-			if(new_version >= ::jconf::inst()->GetMiningForkVersion())
+			coinDescription coinDesc = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(oWork.iPoolId);
+			if(new_version >= coinDesc.GetMiningForkVersion())
 			{
-				miner_algo = ::jconf::inst()->GetMiningAlgo();
+				miner_algo = coinDesc.GetMiningAlgo();
 				hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
 			}
+			else
+			{
+				miner_algo = coinDesc.GetMiningAlgoRoot();
+				hash_fun = func_selector(::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
+			}
+			lastPoolId = oWork.iPoolId;
 			version = new_version;
 		}
 
@@ -627,22 +635,22 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
 
 void minethd::double_work_main()
 {
-	multiway_work_main<2>();
+	multiway_work_main<2u>();
 }
 
 void minethd::triple_work_main()
 {
-	multiway_work_main<3>();
+	multiway_work_main<3u>();
 }
 
 void minethd::quad_work_main()
 {
-	multiway_work_main<4>();
+	multiway_work_main<4u>();
 }
 
 void minethd::penta_work_main()
 {
-	multiway_work_main<5>();
+	multiway_work_main<5u>();
 }
 
 template<size_t N>
@@ -656,7 +664,7 @@ void minethd::prep_multiway_work(uint8_t *bWorkBlob, uint32_t **piNonce)
 	}
 }
 
-template<size_t N>
+template<uint32_t N>
 void minethd::multiway_work_main()
 {
 	if(affinity >= 0) //-1 means no affinity
@@ -689,9 +697,10 @@ void minethd::multiway_work_main()
 	globalStates::inst().iConsumeCnt++;
 
 	// start with root algorithm and switch later if fork version is reached
-	auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+	auto miner_algo = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot();
 	cn_hash_fun_multi hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
 	uint8_t version = 0;
+	size_t lastPoolId = 0;
 
 	while (bQuit == 0)
 	{
@@ -718,13 +727,20 @@ void minethd::multiway_work_main()
 			iNonce = *piNonce[0];
 
 		uint8_t new_version = oWork.getVersion();
-		if(new_version != version)
+		if(new_version != version || oWork.iPoolId != lastPoolId)
 		{
-			if(new_version >= ::jconf::inst()->GetMiningForkVersion())
+			coinDescription coinDesc = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(oWork.iPoolId);
+			if(new_version >= coinDesc.GetMiningForkVersion())
+			{
+				miner_algo = coinDesc.GetMiningAlgo();
+				hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
+			}
+			else
 			{
-				miner_algo = ::jconf::inst()->GetMiningAlgo();
+				miner_algo = coinDesc.GetMiningAlgoRoot();
 				hash_fun_multi = func_multi_selector(N, ::jconf::inst()->HaveHardwareAes(), bNoPrefetch, miner_algo);
 			}
+			lastPoolId = oWork.iPoolId;
 			version = new_version;
 		}
 
@@ -769,4 +785,4 @@ void minethd::multiway_work_main()
 }
 
 } // namespace cpu
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/cpu/minethd.hpp b/xmrstak/backend/cpu/minethd.hpp
index cd78343..85a95d1 100644
--- a/xmrstak/backend/cpu/minethd.hpp
+++ b/xmrstak/backend/cpu/minethd.hpp
@@ -35,7 +35,7 @@ private:
 
 	minethd(miner_work& pWork, size_t iNo, int iMultiway, bool no_prefetch, int64_t affinity);
 
-	template<size_t N>
+	template<uint32_t N>
 	void multiway_work_main();
 
 	template<size_t N>
@@ -65,4 +65,4 @@ private:
 };
 
 } // namespace cpu
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/globalStates.cpp b/xmrstak/backend/globalStates.cpp
index 78823c5..1ec7983 100644
--- a/xmrstak/backend/globalStates.cpp
+++ b/xmrstak/backend/globalStates.cpp
@@ -53,4 +53,4 @@ void globalStates::switch_work(miner_work& pWork, pool_data& dat)
 	iGlobalJobNo++;
 }
 
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/globalStates.hpp b/xmrstak/backend/globalStates.hpp
index 442be69..fafd232 100644
--- a/xmrstak/backend/globalStates.hpp
+++ b/xmrstak/backend/globalStates.hpp
@@ -55,4 +55,4 @@ private:
 	}
 };
 
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/iBackend.hpp b/xmrstak/backend/iBackend.hpp
index 3d2115a..fdc647e 100644
--- a/xmrstak/backend/iBackend.hpp
+++ b/xmrstak/backend/iBackend.hpp
@@ -47,4 +47,4 @@ namespace xmrstak
 		}
 	};
 
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/miner_work.hpp b/xmrstak/backend/miner_work.hpp
index 9e5a4e4..438ec0d 100644
--- a/xmrstak/backend/miner_work.hpp
+++ b/xmrstak/backend/miner_work.hpp
@@ -81,4 +81,4 @@ namespace xmrstak
 		}
 
 	};
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/nvidia/autoAdjust.hpp b/xmrstak/backend/nvidia/autoAdjust.hpp
index be7d1ce..d8bb621 100644
--- a/xmrstak/backend/nvidia/autoAdjust.hpp
+++ b/xmrstak/backend/nvidia/autoAdjust.hpp
@@ -50,7 +50,7 @@ public:
 			ctx.device_blocks = -1;
 			ctx.device_threads = -1;
 
-			// set all evice option those marked as auto (-1) to a valid value
+			// set all device option those marked as auto (-1) to a valid value
 #ifndef _WIN32
 			ctx.device_bfactor = 0;
 			ctx.device_bsleep = 0;
@@ -109,4 +109,4 @@ private:
 };
 
 } // namespace nvidia
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/nvidia/jconf.hpp b/xmrstak/backend/nvidia/jconf.hpp
index 7f60f1d..b4ebaa0 100644
--- a/xmrstak/backend/nvidia/jconf.hpp
+++ b/xmrstak/backend/nvidia/jconf.hpp
@@ -49,4 +49,4 @@ private:
 };
 
 } // namespace nvidia
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/nvidia/minethd.cpp b/xmrstak/backend/nvidia/minethd.cpp
index 804c06a..92f5f78 100644
--- a/xmrstak/backend/nvidia/minethd.cpp
+++ b/xmrstak/backend/nvidia/minethd.cpp
@@ -239,7 +239,7 @@ void minethd::work_main()
 	cpu_ctx = cpu::minethd::minethd_alloc_ctx();
 	
 	// start with root algorithm and switch later if fork version is reached
-	auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+	auto miner_algo = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot();
 	cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
 	
 	uint32_t iNonce;
@@ -247,6 +247,7 @@ void minethd::work_main()
 	globalStates::inst().iConsumeCnt++;
 
 	uint8_t version = 0;
+	size_t lastPoolId = 0;
 
 	while (bQuit == 0)
 	{
@@ -264,13 +265,20 @@ void minethd::work_main()
 			continue;
 		}
 		uint8_t new_version = oWork.getVersion();
-		if(new_version != version)
+		if(new_version != version || oWork.iPoolId != lastPoolId)
 		{
-			if(new_version >= ::jconf::inst()->GetMiningForkVersion())
+			coinDescription coinDesc = ::jconf::inst()->GetCurrentCoinSelection().GetDescription(oWork.iPoolId);
+			if(new_version >= coinDesc.GetMiningForkVersion())
 			{
-				miner_algo = ::jconf::inst()->GetMiningAlgo();
+				miner_algo = coinDesc.GetMiningAlgo();
 				hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
 			}
+			else
+			{
+				miner_algo = coinDesc.GetMiningAlgoRoot();
+				hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
+			}
+			lastPoolId = oWork.iPoolId;
 			version = new_version;
 		}
 
diff --git a/xmrstak/backend/nvidia/minethd.hpp b/xmrstak/backend/nvidia/minethd.hpp
index fcd24fa..89c2944 100644
--- a/xmrstak/backend/nvidia/minethd.hpp
+++ b/xmrstak/backend/nvidia/minethd.hpp
@@ -60,4 +60,4 @@ private:
 };
 
 } // namespace nvidia
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
index 02c157e..f192f01 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
@@ -279,12 +279,15 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx)
 	if(gpuArch < 70)
 		CUDA_CHECK(ctx->device_id, cudaDeviceSetCacheConfig(cudaFuncCachePreferL1));
 
-	size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
+	size_t hashMemSize = std::max(
+		cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()),
+		cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot())
+	);
 
 	size_t wsize = ctx->device_blocks * ctx->device_threads;
 	CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_state, 50 * sizeof(uint32_t) * wsize));
 	size_t ctx_b_size = 4 * sizeof(uint32_t) * wsize;
-	if(cryptonight_heavy == ::jconf::inst()->GetMiningAlgo())
+	if(cryptonight_heavy == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo())
 	{
 		// extent ctx_b to hold the state of idx0
 		ctx_b_size += sizeof(uint32_t) * wsize;
@@ -458,7 +461,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
 #undef XMRSTAK_PP_TOSTRING1
 	std::stringstream ss(archStringList);
 
-	//transform string list sperated with `+` into a vector of integers
+	//transform string list separated with `+` into a vector of integers
 	int tmpArch;
 	while ( ss >> tmpArch )
 		arch.push_back( tmpArch );
@@ -492,7 +495,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
 		}
 	}
 
-	// set all evice option those marked as auto (-1) to a valid value
+	// set all device option those marked as auto (-1) to a valid value
 	if(ctx->device_blocks == -1)
 	{
 		/* good values based of my experience
@@ -576,7 +579,10 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
 		ctx->total_device_memory = totalMemory;
 		ctx->free_device_memory = freeMemory;
 
-		size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
+		size_t hashMemSize = std::max(
+			cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()),
+			cn_select_memory(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgoRoot())
+		);
 
 #ifdef WIN32
 		/* We use in windows bfactor (split slow kernel into smaller parts) to avoid
@@ -606,7 +612,7 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
 		// up to 16kibyte extra memory is used per thread for some kernel (lmem/local memory)
 		// 680bytes are extra meta data memory per hash
 		size_t perThread = hashMemSize + 16192u + 680u;
-		if(cryptonight_heavy == ::jconf::inst()->GetMiningAlgo())
+		if(cryptonight_heavy == ::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo())
 			perThread += 50 * 4; // state double buffer
 		
 		size_t max_intensity = limitedMemory / perThread;
diff --git a/xmrstak/backend/plugin.hpp b/xmrstak/backend/plugin.hpp
index 89cdf97..2610db8 100644
--- a/xmrstak/backend/plugin.hpp
+++ b/xmrstak/backend/plugin.hpp
@@ -109,4 +109,4 @@ struct plugin
  * */
 };
 
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/cli/cli-miner.cpp b/xmrstak/cli/cli-miner.cpp
index c425f04..6118682 100644
--- a/xmrstak/cli/cli-miner.cpp
+++ b/xmrstak/cli/cli-miner.cpp
@@ -80,6 +80,8 @@ void help()
 #ifndef CONF_NO_OPENCL
 	cout<<"  --noAMD                    disable the AMD miner backend"<<endl;
 	cout<<"  --noAMDCache               disable the AMD(OpenCL) cache for precompiled binaries"<<endl;
+	cout<<"  --openCLVendor VENDOR      use OpenCL driver of VENDOR and devices [AMD,NVIDIA]"<<endl;
+	cout<<"                             default: AMD"<<endl;
 	cout<<"  --amd FILE                 AMD backend miner config file"<<endl;
 #endif
 #ifndef CONF_NO_CUDA
@@ -450,6 +452,24 @@ int main(int argc, char *argv[])
 		{
 			params::inst().useAMD = false;
 		}
+		else if(opName.compare("--openCLVendor") == 0)
+		{
+			++i;
+			if( i >=argc )
+			{
+				printer::inst()->print_msg(L0, "No argument for parameter '--openCLVendor' given");
+				win_exit();
+				return 1;
+			}
+			std::string vendor(argv[i]);
+			params::inst().openCLVendor = vendor;
+			if(vendor != "AMD" && vendor != "NVIDIA")
+			{
+				printer::inst()->print_msg(L0, "'--openCLVendor' must be 'AMD' or 'NVIDIA'");
+				win_exit();
+				return 1;
+			}
+		}
 		else if(opName.compare("--noAMDCache") == 0)
 		{
 			params::inst().AMDCache = false;
diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp
index f99698a..3af4028 100644
--- a/xmrstak/jconf.cpp
+++ b/xmrstak/jconf.cpp
@@ -86,30 +86,28 @@ configVal oConfigValues[] = {
 
 constexpr size_t iConfigCnt = (sizeof(oConfigValues)/sizeof(oConfigValues[0]));
 
-struct xmrstak_coin_algo
-{
-	const char* coin_name;
-	xmrstak_algo algo;
-	xmrstak_algo algo_root;
-	uint8_t fork_version;
-	const char* default_pool;
-};
-
-xmrstak_coin_algo coin_algos[] = { 
-	{ "aeon7", cryptonight_aeon, cryptonight_lite, 7u, "mine.aeon-pool.com:5555" },
-	{ "cryptonight", cryptonight, cryptonight, 0u, nullptr },
-	{ "cryptonight_lite", cryptonight_lite, cryptonight_lite, 0u, nullptr },
-	{ "edollar", cryptonight, cryptonight, 0u, nullptr },
-	{ "electroneum", cryptonight, cryptonight, 0u, nullptr },
-	{ "graft", cryptonight, cryptonight, 0u, nullptr },
-	{ "intense", cryptonight, cryptonight, 0u, nullptr },
-	{ "karbo", cryptonight, cryptonight, 0u, nullptr },
-	{ "monero7", cryptonight_monero, cryptonight, 7u, "pool.usxmrpool.com:3333" },
-	{ "stellite", cryptonight_monero, cryptonight, 3u, nullptr },
-	{ "sumokoin", cryptonight_heavy, cryptonight, 3u, nullptr }
+xmrstak::coin_selection coins[] = {
+	// name, userpool, devpool, default_pool_suggestion
+	{ "aeon7",               {cryptonight_aeon, cryptonight_lite, 7u},     {cryptonight_aeon, cryptonight_lite, 7u},     "mine.aeon-pool.com:5555" },
+	{ "bbscoin",             {cryptonight_monero, cryptonight, 3u},        {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+	{ "croat",               {cryptonight_monero, cryptonight, 255u},      {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+	{ "cryptonight",         {cryptonight_monero, cryptonight, 255u},      {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+	{ "cryptonight_heavy",   {cryptonight_heavy, cryptonight_heavy, 0u},   {cryptonight_heavy, cryptonight_heavy, 0u},   nullptr },
+	{ "cryptonight_lite",    {cryptonight_aeon, cryptonight_lite, 255u},   {cryptonight_aeon, cryptonight_lite, 7u},     nullptr },
+	{ "cryptonight_lite_v7", {cryptonight_lite, cryptonight_aeon, 255u},   {cryptonight_aeon, cryptonight_lite, 7u},     nullptr },
+	{ "cryptonight_v7",      {cryptonight_monero, cryptonight_monero, 0u}, {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+	{ "edollar",             {cryptonight_monero, cryptonight, 255u},      {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+	{ "electroneum",         {cryptonight_monero, cryptonight, 255u},      {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+	{ "graft",               {cryptonight_monero, cryptonight, 8u},        {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+	{ "haven",               {cryptonight_heavy, cryptonight, 2u},         {cryptonight_heavy, cryptonight_heavy, 0u},   nullptr },
+	{ "intense",             {cryptonight_monero, cryptonight, 4u},        {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+	{ "karbo",               {cryptonight_monero, cryptonight, 255u},      {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+	{ "monero7",             {cryptonight_monero, cryptonight_monero, 0u}, {cryptonight_monero, cryptonight_monero, 0u}, "pool.usxmrpool.com:3333" },
+	{ "stellite",            {cryptonight_monero, cryptonight, 3u},        {cryptonight_monero, cryptonight_monero, 0u}, nullptr },
+	{ "sumokoin",            {cryptonight_heavy, cryptonight_heavy, 0u},   {cryptonight_heavy, cryptonight_heavy, 0u},   nullptr }
 };
 
-constexpr size_t coin_alogo_size = (sizeof(coin_algos)/sizeof(coin_algos[0]));
+constexpr size_t coin_alogo_size = (sizeof(coins)/sizeof(coins[0]));
 
 inline bool checkType(Type have, Type want)
 {
@@ -313,7 +311,7 @@ void jconf::GetAlgoList(std::string& list)
 	for(size_t i=0; i < coin_alogo_size; i++)
 	{
 		list += "\t- ";
-		list += coin_algos[i].coin_name;
+		list += coins[i].coin_name;
 		list += "\n";
 	}
 }
@@ -331,7 +329,7 @@ bool jconf::IsOnAlgoList(std::string& needle)
 
 	for(size_t i=0; i < coin_alogo_size; i++)
 	{
-		if(needle == coin_algos[i].coin_name)
+		if(needle == coins[i].coin_name)
 			return true;
 	}
 	return false;
@@ -343,10 +341,10 @@ const char* jconf::GetDefaultPool(const char* needle)
 	
 	for(size_t i=0; i < coin_alogo_size; i++)
 	{
-		if(strcmp(needle, coin_algos[i].coin_name) == 0)
+		if(strcmp(needle, coins[i].coin_name) == 0)
 		{
-			if(coin_algos[i].default_pool != nullptr)
-				return coin_algos[i].default_pool;
+			if(coins[i].default_pool != nullptr)
+				return coins[i].default_pool;
 			else
 				return default_example;
 		}
@@ -630,16 +628,14 @@ bool jconf::parse_config(const char* sFilename, const char* sFilenamePools)
 			return false;
 		}
 
-		if(ctmp == coin_algos[i].coin_name)
+		if(ctmp == coins[i].coin_name)
 		{
-			mining_algo = coin_algos[i].algo;
-			mining_algo_root = coin_algos[i].algo_root;
-			mining_fork_version = coin_algos[i].fork_version;
+			currentCoin = coins[i];
 			break;
 		}
 	}
 
-	if(mining_algo == invalid_algo)
+	if(currentCoin.GetDescription(1).GetMiningAlgo() == invalid_algo)
 	{
 		std::string cl;
 		GetAlgoList(cl);
diff --git a/xmrstak/jconf.hpp b/xmrstak/jconf.hpp
index 76b93dd..102b70f 100644
--- a/xmrstak/jconf.hpp
+++ b/xmrstak/jconf.hpp
@@ -1,7 +1,7 @@
 #pragma once
 
-#include "xmrstak/backend/cryptonight.hpp"
 #include "xmrstak/misc/environment.hpp"
+#include "xmrstak/misc/coinDescription.hpp"
 #include "params.hpp"
 
 #include <stdlib.h>
@@ -48,12 +48,8 @@ public:
 
 	bool TlsSecureAlgos();
 
-	inline xmrstak_algo GetMiningAlgo() const { return mining_algo; }
+	inline xmrstak::coin_selection GetCurrentCoinSelection() const { return currentCoin; }
 
-	inline xmrstak_algo GetMiningAlgoRoot() const { return mining_algo_root; }
-
-	inline uint8_t GetMiningForkVersion() const { return mining_fork_version; }
-	
 	std::string GetMiningCoin();
 
 	static void GetAlgoList(std::string& list);
@@ -94,7 +90,5 @@ private:
 	opaque_private* prv;
 
 	bool bHaveAes;
-	xmrstak_algo mining_algo;
-	xmrstak_algo mining_algo_root;
-	uint8_t mining_fork_version;
+	xmrstak::coin_selection currentCoin;
 };
diff --git a/xmrstak/misc/coinDescription.hpp b/xmrstak/misc/coinDescription.hpp
new file mode 100644
index 0000000..73d9a95
--- /dev/null
+++ b/xmrstak/misc/coinDescription.hpp
@@ -0,0 +1,60 @@
+#pragma once
+
+#include "xmrstak/backend/cryptonight.hpp"
+
+#include <stdlib.h>
+#include <string>
+
+
+namespace xmrstak
+{
+	struct coinDescription
+	{
+		xmrstak_algo algo;
+		xmrstak_algo algo_root;
+		uint8_t fork_version;
+
+		coinDescription() = default;
+
+		coinDescription(const xmrstak_algo in_algo, xmrstak_algo in_algo_root, const uint8_t in_fork_version) :
+			algo(in_algo), algo_root(in_algo_root), fork_version(in_fork_version)
+		{}
+
+		inline xmrstak_algo GetMiningAlgo() const { return algo; }
+		inline xmrstak_algo GetMiningAlgoRoot() const { return algo_root; }
+		inline uint8_t GetMiningForkVersion() const { return fork_version; }
+	};
+
+	struct coin_selection
+	{
+		const char* coin_name = nullptr;
+		/* [0] -> user pool
+		 * [1] -> dev pool
+		 */
+		coinDescription pool_coin[2];
+		const char* default_pool = nullptr;
+
+		coin_selection() = default;
+
+		coin_selection(
+			const char* in_coin_name,
+			const coinDescription user_coinDescription,
+			const coinDescription dev_coinDescription,
+			const char* in_default_pool
+		) :
+			coin_name(in_coin_name), default_pool(in_default_pool)
+		{
+			pool_coin[0] = user_coinDescription;
+			pool_coin[1] = dev_coinDescription;
+		}
+
+		/** get coin description for the pool
+		 *
+		 * @param poolId 0 select dev pool, else the user pool is selected
+		 */
+		inline coinDescription GetDescription(size_t poolId) const {
+			coinDescription tmp = (poolId == 0 ? pool_coin[1] : pool_coin[0]);
+			return tmp;
+		}
+	};
+} // namespace xmrstak
diff --git a/xmrstak/misc/configEditor.hpp b/xmrstak/misc/configEditor.hpp
index a840bc4..8a81ad6 100644
--- a/xmrstak/misc/configEditor.hpp
+++ b/xmrstak/misc/configEditor.hpp
@@ -54,4 +54,4 @@ struct configEditor
 
 };
 
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/misc/console.cpp b/xmrstak/misc/console.cpp
index de5eed3..d961b71 100644
--- a/xmrstak/misc/console.cpp
+++ b/xmrstak/misc/console.cpp
@@ -225,7 +225,7 @@ void printer::print_str(const char* str)
 
 //Do a press any key for the windows folk. *insert any key joke here*
 #ifdef _WIN32
-void win_exit(size_t code)
+void win_exit(int code)
 {
 	size_t envSize = 0;
 	getenv_s(&envSize, nullptr, 0, "XMRSTAK_NOWAIT");
@@ -238,7 +238,7 @@ void win_exit(size_t code)
 }
 
 #else
-void win_exit(size_t code) 
+void win_exit(int code)
 { 
 	std::exit(code);
 }
diff --git a/xmrstak/misc/console.hpp b/xmrstak/misc/console.hpp
index cfbeddd..6717631 100644
--- a/xmrstak/misc/console.hpp
+++ b/xmrstak/misc/console.hpp
@@ -49,4 +49,4 @@ private:
 	FILE* logfile;
 };
 
-void win_exit(size_t code = 1);
+void win_exit(int code = 1);
diff --git a/xmrstak/misc/environment.hpp b/xmrstak/misc/environment.hpp
index 99c2db8..b67c858 100644
--- a/xmrstak/misc/environment.hpp
+++ b/xmrstak/misc/environment.hpp
@@ -38,4 +38,4 @@ struct environment
 	params* pParams = nullptr;
 };
 
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp
index 0e1dd9f..8c454eb 100644
--- a/xmrstak/misc/executor.cpp
+++ b/xmrstak/misc/executor.cpp
@@ -421,7 +421,9 @@ void executor::on_miner_result(size_t pool_id, job_result& oResult)
 	{
 		//Ignore errors silently
 		if(pool->is_running() && pool->is_logged_in())
-			pool->cmd_submit(oResult.sJobID, oResult.iNonce, oResult.bResult, backend_name, backend_hashcount, total_hashcount, jconf::inst()->GetMiningAlgo());
+			pool->cmd_submit(oResult.sJobID, oResult.iNonce, oResult.bResult, backend_name, 
+			backend_hashcount, total_hashcount, jconf::inst()->GetCurrentCoinSelection().GetDescription(0).GetMiningAlgo()
+		);
 		return;
 	}
 
@@ -432,7 +434,9 @@ void executor::on_miner_result(size_t pool_id, job_result& oResult)
 	}
 
 	size_t t_start = get_timestamp_ms();
-	bool bResult = pool->cmd_submit(oResult.sJobID, oResult.iNonce, oResult.bResult, backend_name, backend_hashcount, total_hashcount, jconf::inst()->GetMiningAlgo());
+	bool bResult = pool->cmd_submit(oResult.sJobID, oResult.iNonce, oResult.bResult, 
+		backend_name, backend_hashcount, total_hashcount, jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo()
+	);
 	size_t t_len = get_timestamp_ms() - t_start;
 
 	if(t_len > 0xFFFF)
@@ -548,7 +552,7 @@ void executor::ex_main()
 		pools.emplace_back(i+1, params.poolURL.c_str(), params.poolUsername.c_str(), params.poolRigid.c_str(), params.poolPasswd.c_str(), 9.9, false, params.poolUseTls, "", params.nicehashMode);
 	}
 
-	switch(jconf::inst()->GetMiningAlgo())
+	switch(jconf::inst()->GetCurrentCoinSelection().GetDescription(0).GetMiningAlgo())
 	{
 	case cryptonight_heavy:
 		if(dev_tls)
diff --git a/xmrstak/misc/executor.hpp b/xmrstak/misc/executor.hpp
index fbaa265..be5ee6c 100644
--- a/xmrstak/misc/executor.hpp
+++ b/xmrstak/misc/executor.hpp
@@ -23,7 +23,7 @@ namespace cpu
 class minethd;
 
 } // namespace cpu
-} // namepsace xmrstak
+} // namespace xmrstak
 
 class executor
 {
@@ -54,7 +54,7 @@ private:
 
 	inline void set_timestamp() { dev_timestamp = get_timestamp(); };
 
-	// In miliseconds, has to divide a second (1000ms) into an integer number
+	// In milliseconds, has to divide a second (1000ms) into an integer number
 	constexpr static size_t iTickTime = 500;
 
 	// Dev donation time period in seconds. 100 minutes by default.
@@ -64,7 +64,7 @@ private:
 	inline bool is_dev_time()
 	{
 		//Add 2 seconds to compensate for connect
-		constexpr size_t dev_portion = double(iDevDonatePeriod) * fDevDonationLevel + 2;
+		constexpr size_t dev_portion = static_cast<size_t>(double(iDevDonatePeriod) * fDevDonationLevel + 2.);
 
 		if(dev_portion < 12) //No point in bothering with less than 10s
 			return false;
diff --git a/xmrstak/misc/telemetry.cpp b/xmrstak/misc/telemetry.cpp
index 738d287..a3a2a12 100644
--- a/xmrstak/misc/telemetry.cpp
+++ b/xmrstak/misc/telemetry.cpp
@@ -89,8 +89,8 @@ double telemetry::calc_telemetry_data(size_t iLastMilisec, size_t iThread)
 		return nan("");
 
 	double fHashes, fTime;
-	fHashes = iLastestHashCnt - iEarliestHashCnt;
-	fTime = iLastestStamp - iEarliestStamp;
+	fHashes = static_cast<double>(iLastestHashCnt - iEarliestHashCnt);
+	fTime = static_cast<double>(iLastestStamp - iEarliestStamp);
 	fTime /= 1000.0;
 
 	return fHashes / fTime;
@@ -105,4 +105,4 @@ void telemetry::push_perf_value(size_t iThd, uint64_t iHashCount, uint64_t iTime
 	iBucketTop[iThd] = (iTop + 1) & iBucketMask;
 }
 
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/misc/telemetry.hpp b/xmrstak/misc/telemetry.hpp
index b35bbbf..2f84dfa 100644
--- a/xmrstak/misc/telemetry.hpp
+++ b/xmrstak/misc/telemetry.hpp
@@ -21,4 +21,4 @@ private:
 	uint64_t** ppTimestamps;
 };
 
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/misc/utility.cpp b/xmrstak/misc/utility.cpp
index 3b1369a..5177d14 100644
--- a/xmrstak/misc/utility.cpp
+++ b/xmrstak/misc/utility.cpp
@@ -18,4 +18,4 @@ namespace xmrstak
 					}
 				);
 	}
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/misc/utility.hpp b/xmrstak/misc/utility.hpp
index b2e841d..8f2e99f 100644
--- a/xmrstak/misc/utility.hpp
+++ b/xmrstak/misc/utility.hpp
@@ -9,4 +9,4 @@ namespace xmrstak
 	 * @return true if both strings are equal, else false
 	 */
 	bool strcmp_i(const std::string& str1, const std::string& str2);
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp
index d71aeb1..74d1c26 100644
--- a/xmrstak/net/jpsock.cpp
+++ b/xmrstak/net/jpsock.cpp
@@ -229,7 +229,7 @@ void jpsock::jpsock_thread()
 	else
 		disconnect_time = 0;
 
-	std::unique_lock<std::mutex>(job_mutex);
+	std::unique_lock<std::mutex> lck(job_mutex);
 	memset(&oCurrentJob, 0, sizeof(oCurrentJob));
 	bRunning = false;
 }
@@ -439,7 +439,7 @@ bool jpsock::process_pool_job(const opq_json_val* params)
 
 	if(motd != nullptr && motd->IsString() && (motd->GetStringLength() & 0x01) == 0)
 	{
-		std::unique_lock<std::mutex>(motd_mutex);
+		std::unique_lock<std::mutex> lck(motd_mutex);
 		if(motd->GetStringLength() > 0)
 		{
 			pool_motd.resize(motd->GetStringLength()/2 + 1);
@@ -454,7 +454,7 @@ bool jpsock::process_pool_job(const opq_json_val* params)
 
 	executor::inst()->push_event(ex_event(oPoolJob, pool_id));
 
-	std::unique_lock<std::mutex>(job_mutex);
+	std::unique_lock<std::mutex> lck(job_mutex);
 	oCurrentJob = oPoolJob;
 	return true;
 }
@@ -679,13 +679,13 @@ bool jpsock::cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bRes
 
 void jpsock::save_nonce(uint32_t nonce)
 {
-	std::unique_lock<std::mutex>(job_mutex);
+	std::unique_lock<std::mutex> lck(job_mutex);
 	oCurrentJob.iSavedNonce = nonce;
 }
 
 bool jpsock::get_current_job(pool_job& job)
 {
-	std::unique_lock<std::mutex>(job_mutex);
+	std::unique_lock<std::mutex> lck(job_mutex);
 
 	if(oCurrentJob.iWorkLen == 0)
 		return false;
@@ -699,7 +699,7 @@ bool jpsock::get_pool_motd(std::string& strin)
 	if(!ext_motd) 
 		return false;
 
-	std::unique_lock<std::mutex>(motd_mutex);
+	std::unique_lock<std::mutex> lck(motd_mutex);
 	if(pool_motd.size() > 0)
 	{
 		strin.assign(pool_motd);
diff --git a/xmrstak/net/msgstruct.hpp b/xmrstak/net/msgstruct.hpp
index 8c4bdbe..e401f59 100644
--- a/xmrstak/net/msgstruct.hpp
+++ b/xmrstak/net/msgstruct.hpp
@@ -179,7 +179,7 @@ inline size_t get_timestamp()
 	return time_point_cast<seconds>(steady_clock::now()).time_since_epoch().count();
 };
 
-//Get milisecond timestamp
+//Get millisecond timestamp
 inline size_t get_timestamp_ms()
 {
 	using namespace std::chrono;
diff --git a/xmrstak/net/socket.cpp b/xmrstak/net/socket.cpp
index 7c58a8e..9bc608f 100644
--- a/xmrstak/net/socket.cpp
+++ b/xmrstak/net/socket.cpp
@@ -156,7 +156,8 @@ int plain_socket::recv(char* buf, unsigned int len)
 
 bool plain_socket::send(const char* buf)
 {
-	int pos = 0, slen = strlen(buf);
+	size_t pos = 0;
+	size_t slen = strlen(buf);
 	while (pos != slen)
 	{
 		int ret = ::send(hSocket, buf + pos, slen - pos, 0);
diff --git a/xmrstak/params.hpp b/xmrstak/params.hpp
index 3d584b9..6e676cf 100644
--- a/xmrstak/params.hpp
+++ b/xmrstak/params.hpp
@@ -24,6 +24,8 @@ struct params
 	bool AMDCache;
 	bool useNVIDIA;
 	bool useCPU;
+	// user selected OpenCL vendor
+	std::string openCLVendor;
 
 	bool poolUseTls = false;
 	std::string poolURL;
@@ -60,6 +62,7 @@ struct params
 		AMDCache(true),
 		useNVIDIA(true),
 		useCPU(true),
+		openCLVendor("AMD"),
 		configFile("config.txt"),
 		configFilePools("pools.txt"),
 		configFileAMD("amd.txt"),
@@ -69,4 +72,4 @@ struct params
 
 };
 
-} // namepsace xmrstak
+} // namespace xmrstak
diff --git a/xmrstak/pools.tpl b/xmrstak/pools.tpl
index f5afff6..58ca8b5 100644
--- a/xmrstak/pools.tpl
+++ b/xmrstak/pools.tpl
@@ -21,16 +21,27 @@ POOLCONF],
  * Currency to mine. Supported values:
  *
  *    aeon7 (use this for Aeon's new PoW)
- *    cryptonight (try this if your coin is not listed)
- *    cryptonight_lite
+ *    bbscoin (automatic switch with block version 3 to cryptonight_v7)
+ *    croat
  *    edollar
  *    electroneum
  *    graft
+ *    haven
  *    intense
  *    karbo
  *    monero7 (use this for Monero's new PoW)
- *    sumokoin
+ *    sumokoin (automatic switch with block version 3 to cryptonight_heavy)
+ *
+ * Native algorithms which not depends on any block versions:
  *
+ *    # 1MiB scratchpad memory
+ *    cryptonight_lite
+ *    cryptonight_lite_v7
+ *    # 2MiB scratchpad memory
+ *    cryptonight
+ *    cryptonight_v7
+ *    # 4MiB scratchpad memory
+ *    cryptonight_heavy
  */
 
 "currency" : "CURRENCY",
diff --git a/xmrstak/version.cpp b/xmrstak/version.cpp
index 579c27a..f46060c 100644
--- a/xmrstak/version.cpp
+++ b/xmrstak/version.cpp
@@ -18,7 +18,7 @@
 #endif
 
 #define XMR_STAK_NAME "xmr-stak"
-#define XMR_STAK_VERSION "2.4.2"
+#define XMR_STAK_VERSION "2.4.3"
 
 #if defined(_WIN32)
 #define OS_TYPE "win"