2 files changed, 10 insertions, 18 deletions

diff --git a/xmrstak/backend/amd/jconf.cpp b/xmrstak/backend/amd/jconf.cpp
index c2bf1fa..0617aeb 100644
--- a/xmrstak/backend/amd/jconf.cpp
+++ b/xmrstak/backend/amd/jconf.cpp
@@ -48,7 +48,7 @@ using namespace rapidjson;
 /*
  * This enum needs to match index in oConfigValues, otherwise we will get a runtime error
  */
-enum configEnum { iGpuThreadNum, aGpuThreadsConf, iPlatformIdx };
+enum configEnum { aGpuThreadsConf, iPlatformIdx };
 
 struct configVal {
 	configEnum iName;
@@ -58,7 +58,6 @@ struct configVal {
 
 //Same order as in configEnum, as per comment above
 configVal oConfigValues[] = {
-	{ iGpuThreadNum, "gpu_thread_num", kNumberType },
 	{ aGpuThreadsConf, "gpu_threads_conf", kArrayType },
 	{ iPlatformIdx, "platform_index", kNumberType }
 };
@@ -235,14 +234,6 @@ bool jconf::parse_config(const char* sFilename)
 	}
 
 	size_t n_thd = prv->configValues[aGpuThreadsConf]->Size();
-	if(prv->configValues[iGpuThreadNum]->GetUint64() != n_thd)
-	{
-		printer::inst()->print_msg(L0,
-			"Invalid config file. Your GPU config array has %llu members, while you want to use %llu threads.",
-			int_port(n_thd), int_port(prv->configValues[iGpuThreadNum]->GetUint64()));
-		return false;
-	}
-
 	thd_cfg c;
 	for(size_t i=0; i < n_thd; i++)
 	{
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_blake.hpp b/xmrstak/backend/nvidia/nvcc_code/cuda_blake.hpp
index 07ae169..340174c 100644
--- a/xmrstak/backend/nvidia/nvcc_code/cuda_blake.hpp
+++ b/xmrstak/backend/nvidia/nvcc_code/cuda_blake.hpp
@@ -2,7 +2,8 @@
 
 typedef struct {
 	uint32_t h[8], s[4], t[2];
-	int buflen, nullt;
+	uint32_t buflen;
+	int nullt;
 	uint8_t buf[64];
 } blake_state;
 
@@ -50,7 +51,7 @@ __constant__ uint32_t d_blake_cst[16]
 	0xC0AC29B7, 0xC97C50DD, 0x3F84D5B5, 0xB5470917
 };
 
-__device__ void cn_blake_compress(blake_state * __restrict__ S, const uint8_t * __restrict__ block)
+__device__ void cn_blake_compress(blake_state *  S, const uint8_t *  block)
 {
 	uint32_t v[16], m[16], i;
 
@@ -89,12 +90,12 @@ __device__ void cn_blake_compress(blake_state * __restrict__ S, const uint8_t *
 	for (i = 0; i < 8;  ++i) S->h[i] ^= S->s[i % 4];
 }
 
-__device__ void cn_blake_update(blake_state * __restrict__ S, const uint8_t * __restrict__ data, uint64_t datalen)
+__device__ void cn_blake_update(blake_state *  S, const uint8_t *  data, uint64_t datalen)
 {
-	int left = S->buflen >> 3;
-	int fill = 64 - left;
+	uint32_t left = S->buflen >> 3;
+	uint32_t fill = 64 - left;
 
-	if (left && (((datalen >> 3) & 0x3F) >= (unsigned) fill)) 
+	if (left && (((datalen >> 3) & 0x3F) >= fill)) 
 	{
 		memcpy((void *) (S->buf + left), (void *) data, fill);
 		S->t[0] += 512;
@@ -125,7 +126,7 @@ __device__ void cn_blake_update(blake_state * __restrict__ S, const uint8_t * __
 	}
 }
 
-__device__ void cn_blake_final(blake_state * __restrict__ S, uint8_t * __restrict__ digest)
+__device__ void cn_blake_final(blake_state *  S, uint8_t *  digest)
 {
 	const uint8_t padding[] = 
 	{
@@ -177,7 +178,7 @@ __device__ void cn_blake_final(blake_state * __restrict__ S, uint8_t * __restric
 	U32TO8(digest + 28, S->h[7]);
 }
 
-__device__ void cn_blake(const uint8_t * __restrict__ in, uint64_t inlen, uint8_t * __restrict__ out)
+__device__ void cn_blake(const uint8_t *  in, uint64_t inlen, uint8_t *  out)
 {
 	blake_state bs;
 	blake_state *S = (blake_state *)&bs;