Merge 'https://github.com/genesismining/sgminer-gm' to master

Conflicts:
	algorithm.c
	kernel/cryptonight.cl
	util.c

algorithm.c

@@ -1091,11 +1091,12 @@ static cl_int queue_cryptonight_kernel(_clState *clState, dev_blk_ctx *blk, __ma
   cl_ulong le_target = ((cl_ulong)(blk->work->XMRTarget));
 
   //le_target = *(cl_ulong *)(blk->work->device_target + 24);
-  memcpy(clState->cldata, blk->work->data, 76);
+  memcpy(clState->cldata, blk->work->data, blk->work->XMRBlobLen);
 
-  status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 76, clState->cldata , 0, NULL, NULL);
+  status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, blk->work->XMRBlobLen, clState->cldata , 0, NULL, NULL);
 
   CL_SET_ARG(clState->CLbuffer0);
+  CL_SET_ARG(blk->work->XMRBlobLen);
   CL_SET_ARG(clState->Scratchpads);
   CL_SET_ARG(clState->States);
 

algorithm/cryptonight.c

@@ -107,16 +107,22 @@ static void CNKeccakF1600(uint64_t *st)
 	}
 }
 
-void CNKeccak(uint64_t *output, uint64_t *input)
+void CNKeccak(uint64_t *output, uint64_t *input, uint32_t Length)
 {
 	uint64_t st[25];
 
 	// Copy 72 bytes
-	for(int i = 0; i < 9; ++i) st[i] = input[i];
+	//for(int i = 0; i < 9; ++i) st[i] = input[i];
 
-	st[9] = (input[9] & 0x00000000FFFFFFFFUL) | 0x0000000100000000UL;
+	//st[9] = (input[9] & 0x00000000FFFFFFFFUL) | 0x0000000100000000UL;
 
-	for(int i = 10; i < 25; ++i) st[i] = 0x00UL;
+	memcpy(st, input, Length);
+
+	((uint8_t *)st)[Length] = 0x01;
+
+	memset(((uint8_t *)st) + Length + 1, 0x00, 128 - Length - 1);
+
+	for(int i = 16; i < 25; ++i) st[i] = 0x00UL;
 
 	// Last bit of padding
 	st[16] = 0x8000000000000000UL;
@@ -213,13 +219,13 @@ void AESExpandKey256(uint32_t *keybuf)
 	}
 }
 
-void cryptonight(uint8_t *Output, uint8_t *Input)
+void cryptonight(uint8_t *Output, uint8_t *Input, uint32_t Length)
 {
 	CryptonightCtx CNCtx;
 	uint64_t text[16], a[2], b[2];
 	uint32_t ExpandedKey1[64], ExpandedKey2[64];
 
-	CNKeccak(CNCtx.State, Input);
+	CNKeccak(CNCtx.State, Input, Length);
 
 	for(int i = 0; i < 4; ++i) ((uint64_t *)ExpandedKey1)[i] = CNCtx.State[i];
 	for(int i = 0; i < 4; ++i) ((uint64_t *)ExpandedKey2)[i] = CNCtx.State[i + 4];
@@ -334,9 +340,9 @@ void cryptonight_regenhash(struct work *work)
 
 	work->XMRNonce = *nonce;
 
-	memcpy(data, work->data, 19 * 4);
+	memcpy(data, work->data, work->XMRBlobLen);
 
-	cryptonight(ohash, data);
+	cryptonight(ohash, data, work->XMRBlobLen);
 
 	char *tmpdbg = bin2hex(ohash, 32);
 

kernel/cryptonight.cl

@@ -337,7 +337,7 @@ void AESExpandKey256(uint *keybuf)
 #define IDX(x)	((x) * (get_global_size(0)))
 
 __attribute__((reqd_work_group_size(WORKSIZE, 8, 1)))
-__kernel void search(__global ulong *input, __global uint4 *Scratchpad, __global ulong *states)
+__kernel void search(__global ulong *input, uint InputLen, __global uint4 *Scratchpad, __global ulong *states)
 {
 	ulong State[25];
 	uint ExpandedKey1[256];
@@ -364,9 +364,15 @@ __kernel void search(__global ulong *input, __global uint4 *Scratchpad, __global
 	((uint *)State)[9] |= ((get_global_id(0)) & 0xFF) << 24;
 	((uint *)State)[10] &= 0xFFFF0000U;
 	((uint *)State)[10] |= ((get_global_id(0)) & 0xFFFFFF) >> 8;
-	State[9] = (input[9] & 0x00000000FFFFFFFFUL) | 0x0000000100000000UL;
+	State[9] = (input[9] & 0x00000000FFFFFFFFUL); //| 0x0000000100000000UL;
 
-	for(int i = 10; i < 25; ++i) State[i] = 0x00UL;
+	for(int i = 76; i < InputLen; ++i) ((uchar *)State)[i] = ((__global uchar *)input)[i];
+
+	((uchar *)State)[InputLen] = 0x01;
+
+	for(int i = InputLen + 1; i < 128; ++i) ((uchar *)State)[i] = 0x00;
+
+	for(int i = 16; i < 25; ++i) State[i] = 0x00UL;
 
 	// Last bit of padding
 	State[16] = 0x8000000000000000UL;

kernel/ethash-new.cl

@@ -40,7 +40,7 @@ static __constant uint2 const Keccak_f1600_RC[24] = {
 
 #ifdef LEGACY
 #define barrier(x) mem_fence(x)
-#elif WORKSIZE < 64
+#elif WORKSIZE <= 64
 #error "WORKSIZE <= 64 isn't supported by newer AMD drivers and WORKSIZE > 64 is required"
 #endif
 
@@ -212,9 +212,9 @@ __kernel void search(
 
     __local compute_hash_share sharebuf[WORKSIZE / 4];
 #ifdef LEGACY
-    __local buffer[WORKSIZE / 4];
+    __local uint buffer[WORKSIZE / 4];
 #else
-    __local buffer[WORKSIZE];
+    __local uint buffer[WORKSIZE];
 #endif
     __local compute_hash_share * const share = sharebuf + hash_id;
 

miner.h

@@ -1366,7 +1366,8 @@ struct pool {
   //XMR stuff
   char XMRAuthID[64];
   uint32_t XMRTarget;
-  uint8_t XMRBlob[76];
+  uint32_t XMRBlobLen;
+  uint8_t XMRBlob[128];
   pthread_mutex_t XMRGlobalNonceLock;
   uint32_t XMRGlobalNonce;
 
@@ -1528,7 +1529,8 @@ struct work {
 
   /* cryptonight stuff */
   uint32_t XMRTarget;
-  uint8_t XMRBlob[76];
+  uint32_t XMRBlobLen;
+  uint8_t XMRBlob[128];
 
   uint32_t XMRNonce;
 

ocl.c

@@ -1062,7 +1062,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize, algorithm_t *alg
 
   if (algorithm->type == ALGO_CRYPTONIGHT) {
     size_t GlobalThreads;
-    readbufsize = 76UL;
+    readbufsize = 128UL;
 
     set_threads_hashes(1, clState->compute_shaders, &GlobalThreads, 1, &cgpu->intensity, &cgpu->xintensity, &cgpu->rawintensity, &cgpu->algorithm);
 

sgminer.c

@@ -6677,8 +6677,9 @@ static void gen_stratum_work_cn(struct pool *pool, struct work *work)
   work->XMRTarget = pool->XMRTarget;
   //strcpy(work->XMRID, pool->XMRID);
   //work->XMRBlockBlob = strdup(pool->XMRBlockBlob);
-  memcpy(work->XMRBlob, pool->XMRBlob, 76);
-  memcpy(work->data, work->XMRBlob, 76);
+  memcpy(work->XMRBlob, pool->XMRBlob, pool->XMRBlobLen);
+  work->XMRBlobLen = pool->XMRBlobLen;
+  memcpy(work->data, work->XMRBlob, work->XMRBlobLen);
   memset(work->target, 0xFF, 32);
   work->sdiff = (double)0xffffffff / pool->XMRTarget;
   work->work_difficulty = work->sdiff;

util.c

@@ -1960,7 +1960,7 @@ bool parse_notify_cn(struct pool *pool, json_t *val)
   char *job_id = NULL;
   bool clean;
   uint32_t XMRTarget;
-  uint8_t XMRBlob[76];
+  uint8_t XMRBlob[128];
   int ret = true;
 
   /*json_t *arr = json_array_get(val, 0);
@@ -1986,7 +1986,15 @@ bool parse_notify_cn(struct pool *pool, json_t *val)
   }
 
   const char *blobval = json_string_value(blob);
-  if (!hex2bin(XMRBlob, blobval, 76)) {
+
+  if(strlen(blobval) > 254)
+  {
+    applog(LOG_ERR, "Got a massive blob from pool.");
+    ret = false;
+    goto out;
+  }
+
+  if (!hex2bin(XMRBlob, blobval, strlen(blobval) / 2)) {
     ret = false;
     goto out;
   }
@@ -2003,7 +2011,8 @@ bool parse_notify_cn(struct pool *pool, json_t *val)
   pool->swork.job_id = strdup(job_id);
   pool->swork.clean = true;
 
-  memcpy(pool->XMRBlob, XMRBlob, 76);
+  pool->XMRBlobLen = strlen(blobval) >> 1;
+  memcpy(pool->XMRBlob, XMRBlob, pool->XMRBlobLen);
   pool->XMRTarget = XMRTarget;
   pool->swork.diff = (double)0xffffffff / XMRTarget;
   pool->getwork_requested++;