Merge pull request #1279 from fireice-uk/dev

release 2.4.0

4 files changed, 49 insertions, 62 deletions

diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp
index 8d0fd32..2973db4 100644
--- a/xmrstak/backend/amd/amd_gpu/gpu.cpp
+++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp
@@ -16,6 +16,7 @@
 #include "xmrstak/backend/cryptonight.hpp"
 #include "xmrstak/jconf.hpp"
 #include "xmrstak/picosha2/picosha2.hpp"
+#include "xmrstak/params.hpp"
 
 #include <stdio.h>
 #include <string.h>
@@ -393,9 +394,10 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_
 
 	std::string cache_file = get_home() + "/.openclcache/" + hash_hex_str + ".openclbin";
 	std::ifstream clBinFile(cache_file, std::ofstream::in | std::ofstream::binary);
-	if(!clBinFile.good())
+	if(xmrstak::params::inst().AMDCache == false || !clBinFile.good())
 	{
-		printer::inst()->print_msg(L1,"OpenCL device %u - Precompiled code %s not found. Compiling ...",ctx->deviceIdx, cache_file.c_str());
+		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)
 		{
@@ -467,29 +469,31 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_
 		std::vector<char*> all_programs(num_devices);
 		std::vector<std::vector<char>> program_storage;
 
-		int p_id = 0;
-		size_t mem_size = 0;
-		// create memory  structure to query all OpenCL program binaries
-		for(auto & p : all_programs)
+		if(xmrstak::params::inst().AMDCache)
 		{
-			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++;
-		}
+			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, 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((ret = clGetProgramInfo(ctx->Program, 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;
+			}
 
-		std::ofstream file_stream;
-		std::cout<<get_home() + "/.openclcache/" + hash_hex_str + ".openclbin"<<std::endl;
-		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());
+			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
 	{
@@ -936,7 +940,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
 
 	/// @todo only activate if currency is monero
 	int cn_kernel_offset = 0;
-	if(miner_algo == cryptonight_monero && version >= 7)
+	if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon)
 	{
 		cn_kernel_offset = 6;
 	}
@@ -962,7 +966,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
 		return(ERR_OCL_API);
 	}
 
-	if(miner_algo == cryptonight_monero && version >= 7)
+	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)
@@ -1130,7 +1134,7 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, xmrstak_algo miner_algo,
 	size_t tmpNonce = ctx->Nonce;
 	/// @todo only activate if currency is monero
 	int cn_kernel_offset = 0;
-	if(miner_algo == cryptonight_monero && version >= 7)
+	if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon)
 	{
 		cn_kernel_offset = 6;
 	}
diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
index 7a36357..5d4e66c 100644
--- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
+++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
@@ -433,10 +433,7 @@ inline ulong getIdx()
 #endif
 }
 
-inline uint4 mix_and_propagate(__local uint4 xin[8][WORKSIZE])
-{
-	return 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)]
 
 __attribute__((reqd_work_group_size(WORKSIZE, 8, 1)))
 __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states, ulong Threads
diff --git a/xmrstak/backend/amd/jconf.cpp b/xmrstak/backend/amd/jconf.cpp
index 93ba709..0f39ff2 100644
--- a/xmrstak/backend/amd/jconf.cpp
+++ b/xmrstak/backend/amd/jconf.cpp
@@ -234,15 +234,15 @@ bool jconf::parse_config(const char* sFilename)
 
 	if(prv->jsonDoc.HasParseError())
 	{
-		printer::inst()->print_msg(L0, "JSON config parse error(offset %llu): %s",
-			int_port(prv->jsonDoc.GetErrorOffset()), GetParseError_En(prv->jsonDoc.GetParseError()));
+		printer::inst()->print_msg(L0, "JSON config parse error in '%s' (offset %llu): %s",
+			sFilename, int_port(prv->jsonDoc.GetErrorOffset()), GetParseError_En(prv->jsonDoc.GetParseError()));
 		return false;
 	}
 
 
 	if(!prv->jsonDoc.IsObject())
 	{ //This should never happen as we created the root ourselves
-		printer::inst()->print_msg(L0, "Invalid config file. No root?\n");
+		printer::inst()->print_msg(L0, "Invalid config file '%s'. No root?", sFilename);
 		return false;
 	}
 
@@ -258,13 +258,13 @@ bool jconf::parse_config(const char* sFilename)
 
 		if(prv->configValues[i] == nullptr)
 		{
-			printer::inst()->print_msg(L0, "Invalid config file. Missing value \"%s\".", oConfigValues[i].sName);
+			printer::inst()->print_msg(L0, "Invalid config file '%s'. Missing value \"%s\".", sFilename, oConfigValues[i].sName);
 			return false;
 		}
 
 		if(!checkType(prv->configValues[i]->GetType(), oConfigValues[i].iType))
 		{
-			printer::inst()->print_msg(L0, "Invalid config file. Value \"%s\" has unexpected type.", oConfigValues[i].sName);
+			printer::inst()->print_msg(L0, "Invalid config file '%s'. Value \"%s\" has unexpected type.", sFilename, oConfigValues[i].sName);
 			return false;
 		}
 	}
diff --git a/xmrstak/backend/amd/minethd.cpp b/xmrstak/backend/amd/minethd.cpp
index 46a04d5..cab5ad9 100644
--- a/xmrstak/backend/amd/minethd.cpp
+++ b/xmrstak/backend/amd/minethd.cpp
@@ -191,15 +191,10 @@ void minethd::work_main()
 	uint64_t iCount = 0;
 	cryptonight_ctx* cpu_ctx;
 	cpu_ctx = cpu::minethd::minethd_alloc_ctx();
-	auto miner_algo = ::jconf::inst()->GetMiningAlgo();
-	cn_hash_fun hash_fun;
-	if(miner_algo == cryptonight_monero || miner_algo == cryptonight_heavy)
-	{
-		// start with cryptonight and switch later if fork version is reached
-		hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight);
-	}
-	else
-		hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
+	
+	// start with root algorithm and switch later if fork version is reached
+	auto miner_algo = ::jconf::inst()->GetMiningAlgoRoot();
+	cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
 
 	globalStates::inst().iConsumeCnt++;
 
@@ -218,17 +213,18 @@ void minethd::work_main()
 				std::this_thread::sleep_for(std::chrono::milliseconds(100));
 
 			consume_work();
-			uint8_t new_version = oWork.getVersion();
-			if(miner_algo == cryptonight_monero && version < 7 && new_version >= 7)
-			{
-				hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_monero);
-			}
-			else if(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3)
+			continue;
+		}
+
+		uint8_t new_version = oWork.getVersion();
+		if(new_version != version)
+		{
+			if(new_version >= ::jconf::inst()->GetMiningForkVersion())
 			{
-				hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy);
+				miner_algo = ::jconf::inst()->GetMiningAlgo();
+				hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, miner_algo);
 			}
 			version = new_version;
-			continue;
 		}
 
 		uint32_t h_per_round = pGpuCtx->rawIntensity;
@@ -236,7 +232,7 @@ void minethd::work_main()
 
 		assert(sizeof(job_result::sJobID) == sizeof(pool_job::sJobID));
 		uint64_t target = oWork.iTarget;
-		/// \todo add monero hard for version
+		
 		XMRSetJob(pGpuCtx, oWork.bWorkBlob, oWork.iWorkSize, target, miner_algo, version);
 
 		if(oWork.bNiceHash)
@@ -280,16 +276,6 @@ void minethd::work_main()
 		}
 
 		consume_work();
-		uint8_t new_version = oWork.getVersion();
-		if(miner_algo == cryptonight_monero && version < 7 && new_version >= 7)
-		{
-			hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_monero);
-		}
-		else if(miner_algo == cryptonight_heavy && version < 3 && new_version >= 3)
-		{
-			hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, cryptonight_heavy);
-		}
-		version = new_version;
 	}
 }