diff --git a/README.md b/README.md index cf2d787..1708092 100755 --- a/README.md +++ b/README.md @@ -3,6 +3,12 @@ ccminer-cryptonight A modification of Christian Buchner's & Christian H.'s ccminer project by tsiv for Cryptonight mining. +March 31st 2016 +--------------- + +Support for solo mining with bitmonerod v0.9.3.1 or newer. Specify your url +as "daemon+tcp://:/json_rpc" + July 5th 2014 ------------- diff --git a/ccminer.vcxproj b/ccminer.vcxproj index d78f830..3345321 100755 --- a/ccminer.vcxproj +++ b/ccminer.vcxproj @@ -47,7 +47,7 @@ - + @@ -100,7 +100,7 @@ copy "$(CudaToolkitBinDir)\cudart*.dll" "$(OutDir)" true true - compute_20,sm_20;compute_30,sm_30;compute_35,sm_35 + compute_20,sm_20;compute_30,sm_30;compute_35,sm_35;compute_50,sm_50 @@ -131,7 +131,7 @@ copy "$(CudaToolkitBinDir)\cudart*.dll" "$(OutDir)" true true - compute_20,sm_20;compute_30,sm_30;compute_35,sm_35 + compute_20,sm_20;compute_30,sm_30;compute_35,sm_35;compute_50,sm_50 @@ -166,7 +166,7 @@ copy "$(CudaToolkitBinDir)\cudart*.dll" "$(OutDir)" true true - compute_20,sm_20;compute_30,sm_30;compute_35,sm_35 + compute_20,sm_20;compute_30,sm_30;compute_35,sm_35;compute_50,sm_50 @@ -201,7 +201,7 @@ copy "$(CudaToolkitBinDir)\cudart*.dll" "$(OutDir)" true true - compute_20,sm_20;compute_30,sm_30;compute_35,sm_35 + compute_20,sm_20;compute_30,sm_30;compute_35,sm_35;compute_50,sm_50 @@ -287,6 +287,6 @@ copy "$(CudaToolkitBinDir)\cudart*.dll" "$(OutDir)" - + - \ No newline at end of file + diff --git a/cpu-miner.c b/cpu-miner.c old mode 100755 new mode 100644 index 0241d17..66a8f4c --- a/cpu-miner.c +++ b/cpu-miner.c @@ -161,6 +161,7 @@ bool want_longpoll = true; bool have_longpoll = false; bool want_stratum = true; bool have_stratum = false; +bool have_daemon = false; static bool submit_old = false; bool use_syslog = false; static bool opt_background = false; @@ -201,6 +202,7 @@ struct thr_info *thr_info; static int work_thr_id; int longpoll_thr_id = -1; int stratum_thr_id = -1; +int daemon_thr_id = -1; struct work_restart *work_restart = NULL; static struct stratum_ctx stratum; int opt_cn_threads = 8; @@ -496,6 +498,7 @@ bool rpc2_job_decode(const json_t *job, struct work *work) { memset(work->target, 0xff, sizeof(work->target)); work->target[7] = rpc2_target; strncpy(work->job_id, rpc2_job_id, 128); + work->dlen = rpc2_bloblen; } return true; @@ -614,11 +617,12 @@ static void share_result(int result, const char *reason) applog(LOG_DEBUG, "DEBUG: reject reason: %s", reason); } +#define BIG_BUF_LEN 4096 static bool submit_upstream_work(CURL *curl, struct work *work) { char *str = NULL; json_t *val, *res, *reason; - char s[345]; + char s[BIG_BUF_LEN]; int i; bool rc = false; @@ -637,7 +641,7 @@ static bool submit_upstream_work(CURL *curl, struct work *work) if (jsonrpc_2) { noncestr = bin2hex(((const unsigned char*)work->data) + 39, 4); char hash[32]; - cryptonight_hash((void *)hash, (const void *)work->data, 76); + cryptonight_hash((void *)hash, (const void *)work->data, work->dlen); char *hashhex = bin2hex((const unsigned char *)hash, 32); snprintf(s, JSON_BUF_LEN, "{\"method\": \"submit\", \"params\": {\"id\": \"%s\", \"job_id\": \"%s\", \"nonce\": \"%s\", \"result\": \"%s\"}, \"id\":1}\r\n", @@ -671,12 +675,51 @@ static bool submit_upstream_work(CURL *curl, struct work *work) applog(LOG_ERR, "submit_upstream_work stratum_send_line failed"); goto out; } + } else if (have_daemon) { + char *noncestr; + time_t work_time; + pthread_mutex_lock(&g_work_lock); + if (!g_work.xnonce3 || + memcmp(g_work.xnonce3, work->xnonce3, work->xnonce3_len)) { + if (opt_debug) + applog(LOG_DEBUG, "DEBUG: stale work detected, discarding"); + pthread_mutex_unlock(&g_work_lock); + free(work->xnonce3); + work->xnonce3 = NULL; + return true; + } + free(g_work.xnonce3); + g_work.xnonce3 = NULL; + work_time = g_work_time; + pthread_mutex_unlock(&g_work_lock); + noncestr = bin2hex(((const unsigned char*)work->data) + 39, 4); + memcpy(work->xnonce3+78, noncestr, 8); + free(noncestr); + snprintf(s, BIG_BUF_LEN, "{\"method\": \"submitblock\", \"params\": [\"%s\"]}", work->xnonce3); + val = json_rpc_call(curl, rpc_url, NULL, s, false, false, NULL); + if (unlikely(!val)) { + applog(LOG_ERR, "submit_upstream_work json_rpc_call failed"); + goto out; + } + /* tell poller we found a block, get new job */ + tq_push(thr_info[daemon_thr_id].q, (void *)work_time); + res = json_object_get(val, "result"); + if (!res) + { + res = json_object_get(val, "error"); + reason = json_object_get(res, "message"); + } else + reason = NULL; + json_t *status = json_object_get(res, "status"); + share_result(!strcmp(status ? json_string_value(status) : "", "OK"), + reason ? json_string_value(reason) : NULL ); + json_decref(val); } else { /* build JSON-RPC request */ if(jsonrpc_2) { char *noncestr = bin2hex(((const unsigned char*)work->data) + 39, 4); char hash[32]; - cryptonight_hash((void *)hash, (const void *)work->data, 76); + cryptonight_hash((void *)hash, (const void *)work->data, work->dlen); char *hashhex = bin2hex((const unsigned char *)hash, 32); snprintf(s, JSON_BUF_LEN, "{\"method\": \"submit\", \"params\": {\"id\": \"%s\", \"job_id\": \"%s\", \"nonce\": \"%s\", \"result\": \"%s\"}, \"id\":1}\r\n", @@ -925,7 +968,7 @@ static void *workio_thread(void *userdata) return NULL; } - if(!have_stratum) { + if(!have_stratum && !have_daemon) { ok = workio_login(curl); } @@ -1134,11 +1177,13 @@ static void *miner_thread(void *userdata) pthread_mutex_lock(&g_work_lock); if ((*nonceptr) >= end_nonce && !(jsonrpc_2 ? memcmp(work.data, g_work.data, 39) || - memcmp(((uint8_t*) work.data) + 43, ((uint8_t*) g_work.data) + 43, 33) + memcmp(((uint8_t*) work.data) + 43, ((uint8_t*) g_work.data) + 43, work.dlen-43) : memcmp(work.data, g_work.data, 76))) { stratum_gen_work(&stratum, &g_work); } + } else if (have_daemon) { + pthread_mutex_lock(&g_work_lock); } else { /* obtain new work from internal workio thread */ pthread_mutex_lock(&g_work_lock); @@ -1158,10 +1203,14 @@ static void *miner_thread(void *userdata) continue; } } - if (jsonrpc_2 ? memcmp(work.data, g_work.data, 39) || memcmp(((uint8_t*) work.data) + 43, ((uint8_t*) g_work.data) + 43, 33) : memcmp(work.data, g_work.data, 76)) { + if (jsonrpc_2 ? memcmp(work.data, g_work.data, 39) || memcmp(((uint8_t*) work.data) + 43, ((uint8_t*) g_work.data) + 43, work.dlen - 43) : memcmp(work.data, g_work.data, 76)) { memcpy(&work, &g_work, sizeof(struct work)); nonceptr = (uint32_t*) (((char*)work.data) + (jsonrpc_2 ? 39 : 76)); *nonceptr = 0xffffffffU / opt_n_threads * thr_id; + if (work.xnonce3) { + work.xnonce3 = (char *)malloc(work.xnonce3_len); + memcpy(work.xnonce3, g_work.xnonce3, work.xnonce3_len); + } } else ++(*nonceptr); pthread_mutex_unlock(&g_work_lock); @@ -1194,7 +1243,7 @@ static void *miner_thread(void *userdata) gettimeofday(&tv_start, NULL); /* scan nonces for a proof-of-work hash */ - rc = scanhash_cryptonight(thr_id, work.data, work.target, + rc = scanhash_cryptonight(thr_id, work.data, work.dlen, work.target, max_nonce, &hashes_done); // if (opt_benchmark) @@ -1365,6 +1414,123 @@ static void *longpoll_thread(void *userdata) return NULL; } +static const char getblkc[] = "{\"method\": \"getblockcount\"}"; + +#define WALLETLEN 95 + +static void *daemon_thread(void *userdata) { + struct thr_info *mythr = (struct thr_info *)userdata; + CURL *curl = NULL; + uint64_t height, prevheight = 0; + int newblock = 1, delay = 32; + struct timespec ts; + char getblkt[] = "{\"method\": \"getblocktemplate\", \"params\": {\"reserve_size\": 8, \"wallet_address\": " + "\"9xaXMreKDK7bctpHtTE9zUUTgffkRvwZJ7UvyJGAQHkvBFqUYWwhVWWendW6NAdvtB8nn883WQxtU7cpe5eyJiUxLZ741t5\"}}"; + + if (strlen(rpc_user) != WALLETLEN) { + applog(LOG_ERR, "Invalid username / wallet address"); + goto out; + } + + curl = curl_easy_init(); + if (unlikely(!curl)) { + applog(LOG_ERR, "CURL initialization failed"); + goto out; + } + + memcpy(getblkt+80, rpc_user, WALLETLEN); + + applog(LOG_INFO, "Daemon-polling activated for %s", rpc_url); + + while (1) { + json_t *val, *result = NULL, *jheight; + int err; + + /* If we've worked on this job too long, get a new one */ + if (time(0L) > g_work_time + 240) + newblock = 1; + + val = json_rpc_call(curl, rpc_url, NULL, newblock ? getblkt : getblkc, false, false, &err); + if (likely(val)) + result = json_object_get(val, "result"); + if (result) { + time_t qtime; + if (newblock) { + jheight = json_object_get(result, "height"); + if (jheight) { + height = json_integer_value(jheight); + const char *tmpl = json_string_value(json_object_get(result, "blocktemplate_blob")); + const char *hasher = json_string_value(json_object_get(result, "blockhashing_blob")); + uint64_t diff = json_integer_value(json_object_get(result, "difficulty")); + applog(LOG_INFO, "Daemon set diff to %lu on new block", diff); + if (opt_debug) + applog(LOG_DEBUG, "DEBUG: got new work"); + pthread_mutex_lock(&g_work_lock); + g_work.dlen = strlen(hasher)/2; + hex2bin((unsigned char *)g_work.data, hasher, g_work.dlen); + diff = 0xffffffffffffffffUL / diff; + g_work.target[6] = diff & 0xffffffff; + g_work.target[7] = diff >> 32; + free(g_work.xnonce3); + g_work.xnonce3 = strdup(tmpl); + g_work.xnonce3_len = strlen(tmpl)+1; + prevheight = height; + time(&g_work_time); + restart_threads(); + pthread_mutex_unlock(&g_work_lock); + // reduce polling frequency right after + // a new block has been announced. + newblock = 0; + delay = 32; + } + } else { + jheight = json_object_get(result, "count"); + if (jheight) { + height = json_integer_value(jheight); + if (height != prevheight) { + newblock = 1; + json_decref(val); + continue; + } + } + } + json_decref(val); +#ifdef _WIN32 + { struct _timeb now; + _ftime(&now); + ts.tv_nsec = now.millitm * 1000000; + ts.tv_sec = now.time; + } +#else + clock_gettime(CLOCK_REALTIME, &ts); +#endif + ts.tv_sec += delay; +again: + qtime = (time_t)tq_pop(mythr->q, &ts); + if (!qtime) { + /* timed out, adjust delay */ + if (delay > 1) + delay >>= 1; + } else if (qtime == g_work_time) + newblock = 1; + else + goto again; + } else { + if (val) { + result = json_object_get(val, "error"); + applog(LOG_DEBUG, "DEBUG: getblocktemplate failed: %s", json_string_value(json_object_get(result, "message"))); + json_decref(val); + } + restart_threads(); + sleep(opt_fail_pause); + continue; + } + } + +out: + return NULL ; +} + static bool stratum_handle_response(char *buf) { json_t *val, *err_val, *res_val, *id_val; @@ -1620,9 +1786,18 @@ static void parse_arg (int key, char *arg) if (p) { if (strncasecmp(arg, "http://", 7) && strncasecmp(arg, "https://", 8) - && strncasecmp(arg, "stratum+tcp://", 14)) + && strncasecmp(arg, "stratum+tcp://", 14) + && strncasecmp(arg, "daemon+tcp://", 13)) show_usage_and_exit(1); free(rpc_url); + if (!strncasecmp(arg, "daemon", 6)) { + have_daemon = true; + want_longpoll = false; + want_stratum = false; + arg += 6; + arg[0] = 'h'; + arg[2] = 't'; + } rpc_url = strdup(arg); } else { if (!strlen(arg) || *arg == '/') @@ -1996,7 +2171,7 @@ int main(int argc, char *argv[]) if (!work_restart) return 1; - thr_info = (struct thr_info *)calloc(opt_n_threads + 3, sizeof(*thr)); + thr_info = (struct thr_info *)calloc(opt_n_threads + 4, sizeof(*thr)); if (!thr_info) return 1; @@ -2051,6 +2226,21 @@ int main(int argc, char *argv[]) if (have_stratum) tq_push(thr_info[stratum_thr_id].q, strdup(rpc_url)); } + if (have_daemon) { + /* init daemon thread info */ + daemon_thr_id = opt_n_threads + 3; + thr = &thr_info[daemon_thr_id]; + thr->id = daemon_thr_id; + thr->q = tq_new(); + if (!thr->q) + return 1; + + /* start daemon thread */ + if (unlikely(pthread_create(&thr->pth, NULL, daemon_thread, thr))) { + applog(LOG_ERR, "daemon thread create failed"); + return 1; + } + } /* start mining threads */ for (i = 0; i < opt_n_threads; i++) { diff --git a/cryptonight.h b/cryptonight.h index 3977d74..61a441b 100755 --- a/cryptonight.h +++ b/cryptonight.h @@ -1,4 +1,14 @@ +#ifdef __CUDACC__ +# define MY_ALIGN(n) __align__(n) +#elif defined(__GNUC__) +# define MY_ALIGN(n) __attribute__((aligned(n))) +#elif defined(_MSC_VER) +# define MY_ALIGN(n) __declspec(align(n)) +#else +# error "Unknown compiler, need MY_ALIGN definition" +#endif + #define MEMORY (1 << 21) // 2 MiB / 2097152 B #define ITER (1 << 20) // 1048576 #define AES_BLOCK_SIZE 16 @@ -116,7 +126,7 @@ union cn_slow_hash_state { }; }; -struct cryptonight_gpu_ctx { +struct MY_ALIGN(8) cryptonight_gpu_ctx { uint32_t state[50]; uint32_t a[4]; uint32_t b[4]; diff --git a/cryptonight/cryptonight.cu b/cryptonight/cryptonight.cu index 9959a60..72a37d3 100755 --- a/cryptonight/cryptonight.cu +++ b/cryptonight/cryptonight.cu @@ -2,6 +2,7 @@ #include #include #include +#include #include "cuda.h" #include "cuda_runtime.h" @@ -113,13 +114,13 @@ extern void cryptonight_core_cpu_init(int thr_id, int threads); extern void cryptonight_core_cpu_hash(int thr_id, int blocks, int threads, uint32_t *d_long_state, struct cryptonight_gpu_ctx *d_ctx); extern void cryptonight_extra_cpu_setData(int thr_id, const void *data, const void *pTargetIn); -extern void cryptonight_extra_cpu_init(int thr_id); +extern void cryptonight_extra_cpu_init(int thr_id, int dlen); extern void cryptonight_extra_cpu_prepare(int thr_id, int threads, uint32_t startNonce, struct cryptonight_gpu_ctx *d_ctx); extern void cryptonight_extra_cpu_final(int thr_id, int threads, uint32_t startNonce, uint32_t *nonce, struct cryptonight_gpu_ctx *d_ctx); extern "C" void cryptonight_hash(void* output, const void* input, size_t len); -extern "C" int scanhash_cryptonight(int thr_id, uint32_t *pdata, +extern "C" int scanhash_cryptonight(int thr_id, uint32_t *pdata, int dlen, const uint32_t *ptarget, uint32_t max_nonce, unsigned long *hashes_done) { @@ -151,7 +152,7 @@ extern "C" int scanhash_cryptonight(int thr_id, uint32_t *pdata, exit(1); } cryptonight_core_cpu_init(thr_id, throughput); - cryptonight_extra_cpu_init(thr_id); + cryptonight_extra_cpu_init(thr_id, dlen); init[thr_id] = true; } @@ -167,11 +168,11 @@ extern "C" int scanhash_cryptonight(int thr_id, uint32_t *pdata, if (foundNonce < 0xffffffff) { uint32_t vhash64[8]; - uint32_t tempdata[19]; - memcpy(tempdata, pdata, 76); + uint32_t tempdata[32]; uint32_t *tempnonceptr = (uint32_t*)(((char*)tempdata) + 39); + memcpy(tempdata, pdata, dlen); *tempnonceptr = foundNonce; - cryptonight_hash(vhash64, tempdata, 76); + cryptonight_hash(vhash64, tempdata, dlen); if( (vhash64[7] <= Htarg) && fulltest(vhash64, ptarget) ) { diff --git a/cryptonight/cuda_cryptonight_extra.cu b/cryptonight/cuda_cryptonight_extra.cu index 5e6392f..d29226e 100755 --- a/cryptonight/cuda_cryptonight_extra.cu +++ b/cryptonight/cuda_cryptonight_extra.cu @@ -12,6 +12,7 @@ typedef unsigned char BitSequence; typedef unsigned long long DataLength; +static uint32_t d_inlen[8]; static uint32_t *d_input[8]; static uint32_t *d_target[8]; static uint32_t *d_resultNonce[8]; @@ -43,7 +44,7 @@ const uint8_t sub_byte[16][16] = { { 0x8c, 0xa1, 0x89, 0x0d, 0xbf, 0xe6, 0x42, 0x68, 0x41, 0x99, 0x2d, 0x0f, 0xb0, 0x54, 0xbb, 0x16 } }; -__device__ __forceinline__ void cryptonight_aes_set_key( uint32_t *key, const uint32_t *data ) +__device__ void cryptonight_aes_set_key( uint32_t *key, const uint32_t *data ) { int i, j; const int key_base = 8; @@ -74,19 +75,19 @@ __device__ __forceinline__ void cryptonight_aes_set_key( uint32_t *key, const ui } } -__global__ void cryptonight_extra_gpu_prepare(int threads, uint32_t *d_input, uint32_t startNonce, struct cryptonight_gpu_ctx *d_ctx) +__global__ void cryptonight_extra_gpu_prepare(int threads, uint32_t *d_input, int inlen, uint32_t startNonce, struct cryptonight_gpu_ctx *d_ctx) { int thread = (blockDim.x * blockIdx.x + threadIdx.x); if (thread < threads) { struct cryptonight_gpu_ctx ctx; - uint32_t input[19]; + uint32_t input[32]; - MEMCPY4(input, d_input, 19); + memcpy(input, d_input, inlen); *((uint32_t *)(((char *)input) + 39)) = startNonce + thread; - cn_keccak((uint8_t *)input, (uint8_t *)ctx.state); + cn_keccak((uint8_t *)input, inlen, (uint8_t *)ctx.state); cryptonight_aes_set_key(ctx.key1, ctx.state); cryptonight_aes_set_key(ctx.key2, ctx.state+8); XOR_BLOCKS_DST(ctx.state, ctx.state+8, ctx.a); @@ -102,11 +103,11 @@ __global__ void cryptonight_extra_gpu_final(int threads, uint32_t startNonce, ui if (thread < threads) { + MY_ALIGN(8)uint32_t state[50]; int i; uint32_t nonce = startNonce + thread; struct cryptonight_gpu_ctx *ctx = &d_ctx[thread]; uint32_t hash[8]; - uint32_t state[50]; MEMCPY8(state, &ctx->state, 25); cn_keccakf((uint64_t *)state); @@ -131,7 +132,9 @@ __global__ void cryptonight_extra_gpu_final(int threads, uint32_t startNonce, ui int position = -1; bool rc = true; +#if 0 #pragma unroll 8 +#endif for (i = 7; i >= 0; i--) { if (hash[i] > d_target[i]) { if(position < i) { @@ -155,13 +158,14 @@ __global__ void cryptonight_extra_gpu_final(int threads, uint32_t startNonce, ui __host__ void cryptonight_extra_cpu_setData(int thr_id, const void *data, const void *pTargetIn) { - cudaMemcpy(d_input[thr_id], data, 19*sizeof(uint32_t), cudaMemcpyHostToDevice); + cudaMemcpy(d_input[thr_id], data, d_inlen[thr_id], cudaMemcpyHostToDevice); cudaMemcpy(d_target[thr_id], pTargetIn, 8*sizeof(uint32_t), cudaMemcpyHostToDevice); } -__host__ void cryptonight_extra_cpu_init(int thr_id) +__host__ void cryptonight_extra_cpu_init(int thr_id, int inlen) { - cudaMalloc(&d_input[thr_id], 19*sizeof(uint32_t)); + d_inlen[thr_id] = inlen; + cudaMalloc(&d_input[thr_id], inlen); cudaMalloc(&d_target[thr_id], 8*sizeof(uint32_t)); cudaMalloc(&d_resultNonce[thr_id], sizeof(uint32_t)); cudaMemcpyToSymbol(keccakf_rndc, h_keccakf_rndc, sizeof(h_keccakf_rndc), 0, cudaMemcpyHostToDevice); @@ -181,7 +185,7 @@ __host__ void cryptonight_extra_cpu_prepare(int thr_id, int threads, uint32_t st dim3 grid((threads + threadsperblock-1)/threadsperblock); dim3 block(threadsperblock); - cryptonight_extra_gpu_prepare<<>>(threads, d_input[thr_id], startNonce, d_ctx); + cryptonight_extra_gpu_prepare<<>>(threads, d_input[thr_id], d_inlen[thr_id], startNonce, d_ctx); cudaDeviceSynchronize(); } diff --git a/cryptonight/cuda_cryptonight_keccak.cu b/cryptonight/cuda_cryptonight_keccak.cu index 50b6dc0..212b7a7 100755 --- a/cryptonight/cuda_cryptonight_keccak.cu +++ b/cryptonight/cuda_cryptonight_keccak.cu @@ -36,7 +36,7 @@ const uint64_t h_keccakf_rndc[24] = #define rotl64_2(x, y) rotl64_1(((x) >> 32) | ((x) << 32), (y)) #define bitselect(a, b, c) ((a) ^ ((c) & ((b) ^ (a)))) -__device__ __forceinline__ void cn_keccakf(uint64_t *s) +__device__ void cn_keccakf(uint64_t *s) { uint8_t i; @@ -90,13 +90,14 @@ __device__ __forceinline__ void cn_keccakf(uint64_t *s) } } -__device__ __forceinline__ void cn_keccak(const uint8_t * __restrict__ in, uint8_t * __restrict__ md) +__device__ void cn_keccak(const uint8_t * __restrict__ in, int inlen, uint8_t * __restrict__ md) { - uint64_t st[25]; - - MEMCPY4(st, in, 19); - MEMSET8(&st[10], 0x00, 15); - st[9] = (st[9] & 0x00000000FFFFFFFFULL) | 0x0000000100000000ULL; + __align__(8) uint64_t st[25]; + uint8_t *sp = (uint8_t *)st; + + MEMSET8(&st[9], 0x00, 16); + memcpy(sp, in, inlen); + sp[inlen++] = 1; st[16] = 0x8000000000000000ULL; cn_keccakf(st); diff --git a/miner.h b/miner.h index 60b77ae..2c8caaa 100755 --- a/miner.h +++ b/miner.h @@ -194,7 +194,7 @@ void sha256_init_8way(uint32_t *state); void sha256_transform_8way(uint32_t *state, const uint32_t *block, int swap); #endif -extern int scanhash_cryptonight(int thr_id, uint32_t *pdata, +extern int scanhash_cryptonight(int thr_id, uint32_t *pdata, int dlen, const uint32_t *ptarget, uint32_t max_nonce, unsigned long *hashes_done); @@ -253,6 +253,9 @@ struct work { char job_id[128]; size_t xnonce2_len; unsigned char xnonce2[32]; + size_t xnonce3_len; + char *xnonce3; + int dlen; }; struct stratum_job {