Miner Patch for 127 Transaction Bug

The lengths of the buffers for the XMR blog were extended; several
variables to keep track of its size were added; GPU and CPU hashing code
was amended, to take variable-length inputs (from 76 bytes to 128
bytes).

Please refer to fireice_uk's terrific Reddit post on the subject:
https://www.reddit.com/r/Monero/comments/5ptm2y/security_avisory_common_exploit_in_monero_miners/

algorithm.c

@@ -1097,11 +1097,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;
@@ -178,13 +184,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];
@@ -299,9 +305,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] &= 0xFF000000U;
 	((uint *)State)[10] |= ((get_global_id(0) >> 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;

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;
 

sgminer.c

@@ -6668,8 +6668,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

@@ -1972,7 +1972,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);
@@ -1998,7 +1998,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;
   }
@@ -2015,7 +2023,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++;