From 608c9885cbaa89b5f2846e8d1661dda02d2ba8ae Mon Sep 17 00:00:00 2001 From: OhGodAGirl Date: Thu, 9 Feb 2017 00:05:42 +0100 Subject: [PATCH] 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 | 5 +++-- algorithm/cryptonight.c | 22 ++++++++++++++-------- kernel/cryptonight.cl | 12 +++++++++--- miner.h | 6 ++++-- sgminer.c | 5 +++-- util.c | 15 ++++++++++++--- 6 files changed, 45 insertions(+), 20 deletions(-) diff --git a/algorithm.c b/algorithm.c index c2610fbc2..66a54fecd 100644 --- a/algorithm.c +++ b/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); diff --git a/algorithm/cryptonight.c b/algorithm/cryptonight.c index da764b953..ebc9c8a51 100644 --- a/algorithm/cryptonight.c +++ b/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); diff --git a/kernel/cryptonight.cl b/kernel/cryptonight.cl index 2c87dac0e..11d7e26e9 100644 --- a/kernel/cryptonight.cl +++ b/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; diff --git a/miner.h b/miner.h index c59ed7b00..fa63e3eed 100644 --- a/miner.h +++ b/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; diff --git a/sgminer.c b/sgminer.c index 8ef550eda..c61d38534 100644 --- a/sgminer.c +++ b/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; diff --git a/util.c b/util.c index 70cb1902e..6fa921628 100644 --- a/util.c +++ b/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++;