From 806c3e8691215d15fa1dc3f56ba5fcbb6bc21291 Mon Sep 17 00:00:00 2001 From: Tanguy Pruvot Date: Thu, 4 Sep 2014 08:55:00 +0200 Subject: [PATCH] enhance double scan checks --- blake32.cu | 19 +++++++---- cpu-miner.c | 93 ++++++++++++++++++++++++++++++++--------------------- hashlog.cpp | 89 +++++++++++++++++++++++++++++++++++++++----------- miner.h | 1 + util.c | 2 +- 5 files changed, 141 insertions(+), 63 deletions(-) diff --git a/blake32.cu b/blake32.cu index 2b63ccf..ccba68a 100644 --- a/blake32.cu +++ b/blake32.cu @@ -29,8 +29,9 @@ extern "C" void blake32hash(void *output, const void *input) #include "cuda_helper.h" // in cpu-miner.c +extern bool opt_n_threads; extern bool opt_benchmark; -extern bool opt_debug; +//extern bool opt_debug; extern int device_map[8]; extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id); @@ -279,7 +280,9 @@ extern "C" int scanhash_blake32(int thr_id, uint32_t *pdata, const uint32_t *pta ((uint32_t*)ptarget)[7] = 0x00000f; if (!init[thr_id]) { - CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); + if (opt_n_threads > 1) { + CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id])); + } CUDA_SAFE_CALL(cudaMallocHost(&h_resNounce[thr_id], sizeof(uint32_t))); CUDA_SAFE_CALL(cudaMalloc(&d_resNounce[thr_id], sizeof(uint32_t))); init[thr_id] = true; @@ -288,11 +291,6 @@ extern "C" int scanhash_blake32(int thr_id, uint32_t *pdata, const uint32_t *pta if (throughput < (TPB * 2048)) applog(LOG_WARNING, "throughput=%u, start=%x, max=%x", throughput, first_nonce, max_nonce); - if (max_nonce < first_nonce) { - applog(LOG_ERR, "start=%x > end=%x !", first_nonce, max_nonce); - return 0; - } - blake256_cpu_setBlock_80(pdata, (void*)ptarget); do { @@ -340,5 +338,12 @@ extern "C" int scanhash_blake32(int thr_id, uint32_t *pdata, const uint32_t *pta exit_scan: *hashes_done = pdata[19] - first_nonce + 1; + // reset the device to allow multiple instances + if (opt_n_threads == 1) { + CUDA_SAFE_CALL(cudaDeviceReset()); + init[thr_id] = false; + } + // wait proper end of all threads + cudaDeviceSynchronize(); return rc; } diff --git a/cpu-miner.c b/cpu-miner.c index cb4d365..3f772c9 100644 --- a/cpu-miner.c +++ b/cpu-miner.c @@ -187,7 +187,7 @@ static int opt_scantime = 5; static json_t *opt_config; static const bool opt_time = true; static sha256_algos opt_algo = ALGO_HEAVY; -static int opt_n_threads = 0; +int opt_n_threads = 0; static double opt_difficulty = 1; // CH bool opt_trust_pool = false; uint16_t opt_vote = 9999; @@ -411,11 +411,11 @@ err_out: return false; } -static void share_result(int result, const char *reason) +static int share_result(int result, const char *reason) { char s[345]; double hashrate; - int i; + int i, ret = 0; hashrate = 0.; pthread_mutex_lock(&stats_lock); @@ -434,9 +434,15 @@ static void share_result(int result, const char *reason) (result ? CL_GRN "yay!!!" : CL_RED "booooo") : (result ? "(yay!!!)" : "(booooo)")); - if (reason) { + if (reason && !opt_quiet) { applog(LOG_WARNING, "reject reason: %s", reason); + if (strncmp(reason, "low difficulty share", 20) == 0) { + opt_difficulty = (opt_difficulty * 2.0) / 3.0; + applog(LOG_WARNING, "factor reduced to : %0.2f", opt_difficulty); + return 0; + } } + return 1; } static bool submit_upstream_work(CURL *curl, struct work *work) @@ -472,8 +478,10 @@ static bool submit_upstream_work(CURL *curl, struct work *work) sent = hashlog_already_submittted(work->job_id, nonce); if (sent > 0) { sent = (uint32_t) time(NULL) - sent; - if (!opt_quiet) + if (!opt_quiet) { applog(LOG_WARNING, "skip submit, nonce %s was already sent %u seconds ago", noncestr, sent); + hashlog_dump_job(work->job_id); + } rc = true; goto out; } @@ -481,11 +489,11 @@ static bool submit_upstream_work(CURL *curl, struct work *work) if (opt_algo == ALGO_HEAVY) { sprintf(s, "{\"method\": \"mining.submit\", \"params\": [\"%s\", \"%s\", \"%s\", \"%s\", \"%s\", \"%s\"], \"id\":4}", - rpc_user, work->job_id, xnonce2str, ntimestr, noncestr, nvotestr); + rpc_user, work->job_id + 8, xnonce2str, ntimestr, noncestr, nvotestr); } else { sprintf(s, "{\"method\": \"mining.submit\", \"params\": [\"%s\", \"%s\", \"%s\", \"%s\", \"%s\"], \"id\":4}", - rpc_user, work->job_id, xnonce2str, ntimestr, noncestr); + rpc_user, work->job_id + 8, xnonce2str, ntimestr, noncestr); } free(ntimestr); free(noncestr); @@ -528,7 +536,8 @@ static bool submit_upstream_work(CURL *curl, struct work *work) res = json_object_get(val, "result"); reason = json_object_get(val, "reject-reason"); - share_result(json_is_true(res), reason ? json_string_value(reason) : NULL); + if (!share_result(json_is_true(res), reason ? json_string_value(reason) : NULL)) + hashlog_purge_job(work->job_id); json_decref(val); } @@ -768,7 +777,9 @@ static void stratum_gen_work(struct stratum_ctx *sctx, struct work *work) pthread_mutex_lock(&sctx->work_lock); - strcpy(work->job_id, sctx->job.job_id); + // store the job ntime as high part of jobid + snprintf(work->job_id, sizeof(work->job_id), "%07x %s", + be32dec(sctx->job.ntime) & 0xfffffff, sctx->job.job_id); work->xnonce2_len = sctx->xnonce2_size; memcpy(work->xnonce2, sctx->job.xnonce2, sctx->xnonce2_size); @@ -884,8 +895,6 @@ static void *miner_thread(void *userdata) int wcmplen = 76; uint32_t *nonceptr = (uint32_t*) (((char*)work.data) + wcmplen); - applog(LOG_WARNING, "job %s %08x", g_work.job_id, (*nonceptr)); - if (have_stratum) { while (time(NULL) >= (g_work_time + opt_scantime) && !work_done) usleep(500*1000); @@ -909,29 +918,34 @@ static void *miner_thread(void *userdata) g_work_time = time(NULL); } } - if (memcmp(work.data, g_work.data, wcmplen)) { - /* - applog(LOG_NOTICE, "job %s %08x work change", g_work.job_id, (*nonceptr)); - for (int n=0; n 0 && range.scanned[1] > 0) { + } else if (range.scanned[0] > 0 && range.scanned[1] > 0 && range.scanned[1] < 0xFFFFFFF0UL) { /* continue scan the end */ start_nonce = range.scanned[1] + 1; - applog(LOG_WARNING, "scan the next part %x + 1", range.scanned[1]); - } else if (range.scanned[0] > 1) { - /* dont scan the beginning... make loops */ - //end_nonce = range.scanned[0] - 1; - //applog(LOG_WARNING, "scan the missing part 0 -> %x", end_nonce); + //applog(LOG_DEBUG, "scan the next part %x + 1 (%x-%x)", range.scanned[1], range.scanned[0], range.scanned[1]); } - if (start_nonce == work.scanned_from) { - /* to prevent stales, if last was in the same range */ - applog(LOG_ERR, "detected a staled job!"); - //(*nonceptr) = end_nonce + 1; - //work_done = true; - //continue; - start_nonce = range.scanned[1] + 1; + + stall = (start_nonce == work.scanned_from && end_nonce == work.scanned_to); + stall |= (start_nonce == work.scanned_from && start_nonce == range.scanned[1] + 1); + stall |= (start_nonce > range.scanned[0] && start_nonce < range.scanned[1]); + + if (stall) { + if (opt_algo) + applog(LOG_DEBUG, "job done, wait for a new one..."); + work_restart[thr_id].restart = 1; + hashlog_purge_old(); + // wait a bit for a new job... + usleep(1500*1000); + (*nonceptr) = end_nonce + 1; + work_done = true; + continue; } } } diff --git a/hashlog.cpp b/hashlog.cpp index 645fc88..b069d26 100644 --- a/hashlog.cpp +++ b/hashlog.cpp @@ -9,10 +9,12 @@ #define MK_HI64(u32) (0x100000000ULL * u32) struct hashlog_data { - uint32_t ntime; + uint32_t tm_sent; uint32_t scanned_from; uint32_t scanned_to; uint32_t last_from; + uint32_t tm_add; + uint32_t tm_upd; }; static std::map tlastshares; @@ -41,7 +43,7 @@ extern "C" uint32_t hashlog_already_submittted(char* jobid, uint32_t nonce) ret = hashlog_get_last_sent(jobid); } else if (tlastshares.find(key) != tlastshares.end()) { hashlog_data data = tlastshares[key]; - ret = data.ntime; + ret = data.tm_sent; } return ret; } @@ -56,7 +58,9 @@ extern "C" void hashlog_remember_submit(char* jobid, uint32_t nonce) struct hashlog_data data; data = tlastshares[keyall]; - data.ntime = (uint32_t) time(NULL); + data.tm_upd = data.tm_sent = (uint32_t) time(NULL); + if (data.tm_add == 0) + data.tm_add = data.tm_upd; tlastshares[key] = data; } @@ -67,24 +71,38 @@ extern "C" void hashlog_remember_scan_range(char* jobid, uint32_t scanned_from, { uint64_t njobid = hextouint(jobid); uint64_t keyall = (njobid << 32); + uint64_t range = hashlog_get_scan_range(jobid); struct hashlog_data data; // global scan range of a job data = tlastshares[keyall]; - if (hashlog_get_scan_range(jobid) == 0) { + if (range == 0) { memset(&data, 0, sizeof(data)); + } else { + // get min and max from all sent records + data.scanned_from = LO_DWORD(range); + data.scanned_to = HI_DWORD(range); } - if (data.scanned_from == 0 || scanned_to == (data.scanned_from - 1)) - data.scanned_from = scanned_from ? scanned_from : 1; // min 1 - if (data.scanned_to == 0 || scanned_from == data.scanned_to + 1) - data.scanned_to = scanned_to; + if (data.tm_add == 0) + data.tm_add = (uint32_t) time(NULL); data.last_from = scanned_from; + if (scanned_from < scanned_to) { + if (data.scanned_from == 0) + data.scanned_from = scanned_from ? scanned_from : 1; // min 1 + else if (scanned_from < data.scanned_from) // || scanned_to == (data.scanned_from - 1) + data.scanned_from = scanned_from; + if (data.scanned_to == 0 || scanned_from == data.scanned_to + 1) + data.scanned_to = scanned_to; + } + + data.tm_upd = (uint32_t) time(NULL); + tlastshares[keyall] = data; - applog(LOG_BLUE, "job %s range : %x %x -> %x %x (%x)", jobid, - scanned_from, scanned_to, data.scanned_from, data.scanned_to, data.ntime);/* */ + applog(LOG_BLUE, "job %s range : %x %x -> %x %x", jobid, + scanned_from, scanned_to, data.scanned_from, data.scanned_to);/* */ } /** @@ -96,15 +114,21 @@ extern "C" uint64_t hashlog_get_scan_range(char* jobid) uint64_t ret = 0; uint64_t njobid = hextouint(jobid); uint64_t keypfx = (njobid << 32); + struct hashlog_data data; + data.scanned_from = 0; + data.scanned_to = 0; std::map::iterator i = tlastshares.begin(); while (i != tlastshares.end()) { - if ((keypfx & i->first) == keypfx) { - hashlog_data data = i->second; - ret = data.scanned_from; - ret += MK_HI64(data.scanned_to); + if ((keypfx & i->first) == keypfx && i->second.scanned_to > ret) { + if (i->second.scanned_to > data.scanned_to) + data.scanned_to = i->second.scanned_to; + if (i->second.scanned_from < data.scanned_from || data.scanned_from == 0) + data.scanned_from = i->second.scanned_from; } i++; } + ret = data.scanned_from; + ret += MK_HI64(data.scanned_to); return ret; } @@ -119,7 +143,7 @@ extern "C" uint32_t hashlog_get_last_sent(char* jobid) uint64_t keypfx = (njobid << 32); std::map::iterator i = tlastshares.begin(); while (i != tlastshares.end()) { - if ((keypfx & i->first) == keypfx && i->second.ntime > 0) { + if ((keypfx & i->first) == keypfx && i->second.tm_sent > 0) { nonce = LO_DWORD(i->first); } i++; @@ -128,18 +152,25 @@ extern "C" uint32_t hashlog_get_last_sent(char* jobid) } /** - * Remove entries of a job... not used yet + * Remove entries of a job... */ extern "C" void hashlog_purge_job(char* jobid) { + int deleted = 0; uint64_t njobid = hextouint(jobid); uint64_t keypfx = (njobid << 32); + uint32_t sz = tlastshares.size(); std::map::iterator i = tlastshares.begin(); while (i != tlastshares.end()) { - if ((keypfx & i->first) == keypfx) + if ((keypfx & i->first) == keypfx) { + deleted++; tlastshares.erase(i); + } i++; } + if (opt_debug && deleted) { + applog(LOG_DEBUG, "hashlog: purge job %s, del %d/%d", jobid, deleted, sz); + } } /** @@ -152,7 +183,7 @@ extern "C" void hashlog_purge_old(void) uint32_t sz = tlastshares.size(); std::map::iterator i = tlastshares.begin(); while (i != tlastshares.end()) { - if ((now - i->second.ntime) > LOG_PURGE_TIMEOUT) { + if ((now - i->second.tm_sent) > LOG_PURGE_TIMEOUT) { deleted++; tlastshares.erase(i); } @@ -170,3 +201,25 @@ extern "C" void hashlog_purge_all(void) { tlastshares.clear(); } + + +/** + * Can be used to debug... + */ +extern "C" void hashlog_dump_job(char* jobid) +{ + int deleted = 0; + uint64_t njobid = hextouint(jobid); + uint64_t keypfx = (njobid << 32); + uint32_t sz = tlastshares.size(); + std::map::iterator i = tlastshares.begin(); + while (i != tlastshares.end()) { + if ((keypfx & i->first) == keypfx) { + applog(LOG_BLUE, "job %s range : %x %x %s added %x upd %x", jobid, + i->second.scanned_from, i->second.scanned_to, + i->second.tm_sent ? "sent" : "", + i->second.tm_add, i->second.tm_upd);/* */ + } + i++; + } +} \ No newline at end of file diff --git a/miner.h b/miner.h index 5f9e8ac..098b6d5 100644 --- a/miner.h +++ b/miner.h @@ -399,6 +399,7 @@ uint64_t hashlog_get_scan_range(char* jobid); void hashlog_purge_old(void); void hashlog_purge_job(char* jobid); void hashlog_purge_all(void); +void hashlog_dump_job(char* jobid); struct thread_q; diff --git a/util.c b/util.c index 5567459..9afc308 100644 --- a/util.c +++ b/util.c @@ -115,7 +115,7 @@ void applog(int prio, const char *fmt, ...) case LOG_WARNING: color = CL_YLW; break; case LOG_NOTICE: color = CL_WHT; break; case LOG_INFO: color = ""; break; - case LOG_DEBUG: color = ""; break; + case LOG_DEBUG: color = CL_SIL; break; case LOG_BLUE: prio = LOG_NOTICE;