enhance double scan checks

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;
 }

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<wcmplen; n+=8) {
-				if (memcmp(work.data + n, g_work.data + n, 8)) {
-					applog(LOG_ERR, "diff detected at offset %d", n);
-					applog_hash(work.data + n);
-					applog_hash(g_work.data + n);
+		if (memcmp(work.data, g_work.data, 72)) { // wcmplen)) {
+			if (opt_debug) {
+				applog(LOG_DEBUG, "job %s %08x work updated", g_work.job_id, (*nonceptr));
+				for (int n=0; n<wcmplen; n+=8) {
+					if (memcmp(work.data + n, g_work.data + n, 8)) {
+						applog(LOG_DEBUG, "diff detected at offset %d", n);
+						applog_hash((uint8_t*) work.data + n);
+						applog_hash((uint8_t*) g_work.data + n);
+					}
 				}
 			}
-	*/
 			memcpy(&work, &g_work, sizeof(struct work));
 			(*nonceptr) = (0xffffffffUL / opt_n_threads) * thr_id; // 0 if single thr
-	/*	} else if (memcmp(work.target, g_work.target, sizeof(work.target))) {
-			applog(LOG_NOTICE, "job %s %08x target change", g_work.job_id, (*nonceptr));
+		} else if (memcmp(work.target, g_work.target, sizeof(work.target))) {
+			if (opt_debug) {
+				applog(LOG_DEBUG, "job %s %08x target change", g_work.job_id, (*nonceptr));
+				applog_hash((uint8_t*) work.target);
+				applog_hash((uint8_t*) g_work.target);
+			}
 			memcpy(work.target, g_work.target, sizeof(work.target));
 			(*nonceptr) = (0xffffffffUL / opt_n_threads) * thr_id; // 0 if single thr
-	*/
 		} else
 			(*nonceptr)++; //??
 		pthread_mutex_unlock(&g_work_lock);
 		work_restart[thr_id].restart = 0;
 
+		applog(LOG_WARNING, "job %s %08x", g_work.job_id, (*nonceptr));
+
 		/* adjust max_nonce to meet target scan time */
 		if (have_stratum)
 			max64 = LP_SCANTIME;
@@ -959,7 +973,7 @@ static void *miner_thread(void *userdata)
 		start_nonce = *nonceptr;
 
 		/* do not recompute something already scanned */
-		if (opt_algo == ALGO_BLAKE) {
+		if (opt_algo == ALGO_BLAKE && opt_n_threads == 1) {
 			union {
 				uint64_t data;
 				uint32_t scanned[2];
@@ -967,24 +981,29 @@ static void *miner_thread(void *userdata)
 
 			range.data = hashlog_get_scan_range(work.job_id);
 			if (range.data) {
+				bool stall = false;
 				if (range.scanned[0] == 1 && range.scanned[1] == 0xFFFFFFFFUL) {
 					applog(LOG_WARNING, "detected a rescan of fully scanned job!");
-				} else if (range.scanned[0] > 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;
 				}
 			}
 		}

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<uint64_t, hashlog_data> 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<uint64_t, hashlog_data>::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<uint64_t, hashlog_data>::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<uint64_t, hashlog_data>::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<uint64_t, hashlog_data>::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<uint64_t, hashlog_data>::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++;
+	}
+}

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;
 

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;