From ad90269508a9ab85e8b2f5853387720c856ff82b Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Thu, 23 Aug 2012 16:33:24 +1000
Subject: [PATCH 01/17] Minimise the number of getwork threads we generate.

---
 cgminer.c | 14 ++++++++++++++
 1 file changed, 14 insertions(+)

diff --git a/cgminer.c b/cgminer.c
index 5b15e859..508746b3 100644
--- a/cgminer.c
+++ b/cgminer.c
@@ -187,6 +187,7 @@ int hw_errors;
 int total_accepted, total_rejected, total_diff1;
 int total_getworks, total_stale, total_discarded;
 static int total_queued, staged_rollable;
+static int queued_getworks;
 unsigned int new_blocks;
 static unsigned int work_block;
 unsigned int found_blocks;
@@ -2460,6 +2461,9 @@ out:
 	workio_cmd_free(wc);
 	if (ce)
 		push_curl_entry(ce, pool);
+	mutex_lock(&control_lock);
+	queued_getworks--;
+	mutex_unlock(&control_lock);
 	return NULL;
 }
 
@@ -3922,7 +3926,17 @@ static void pool_resus(struct pool *pool)
 static bool queue_request(struct thr_info *thr, bool needed)
 {
 	struct workio_cmd *wc;
+	bool doq = true;
 
+	mutex_lock(&control_lock);
+	if (queued_getworks > (mining_threads + opt_queue) * total_pools)
+		doq = false;
+	else
+		queued_getworks++;
+	mutex_unlock(&control_lock);
+	if (!doq)
+		return true;
+	
 	/* fill out work request message */
 	wc = calloc(1, sizeof(*wc));
 	if (unlikely(!wc)) {

From 4ca288e8204ac8e0fec47af768e3bce37e733ebf Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Thu, 23 Aug 2012 16:55:14 +1000
Subject: [PATCH 02/17] Limit queued_getworks to double the expected queued
 maximum rather than factoring in number of pools.

---
 cgminer.c | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cgminer.c b/cgminer.c
index 508746b3..29362c9b 100644
--- a/cgminer.c
+++ b/cgminer.c
@@ -3929,7 +3929,7 @@ static bool queue_request(struct thr_info *thr, bool needed)
 	bool doq = true;
 
 	mutex_lock(&control_lock);
-	if (queued_getworks > (mining_threads + opt_queue) * total_pools)
+	if (queued_getworks > (mining_threads + opt_queue) * 2)
 		doq = false;
 	else
 		queued_getworks++;

From ec522bdfdbefb0380fa8d8148dee09c03cf2fed6 Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Thu, 23 Aug 2012 17:02:24 +1000
Subject: [PATCH 03/17] Bump to v2.7.3 hotfix.

---
 NEWS         | 5 +++++
 configure.ac | 2 +-
 2 files changed, 6 insertions(+), 1 deletion(-)

diff --git a/NEWS b/NEWS
index dbd8b9cd..adf0a2e3 100644
--- a/NEWS
+++ b/NEWS
@@ -1,3 +1,8 @@
+Version 2.7.3 - August 22, 2012
+
+- Minimise the number of getwork threads we generate.
+
+
 Version 2.7.2 - August 22, 2012
 
 - Pick worksize 256 with Cypress if none is specified.
diff --git a/configure.ac b/configure.ac
index ae5d677f..7ee4885d 100644
--- a/configure.ac
+++ b/configure.ac
@@ -2,7 +2,7 @@
 ##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##
 m4_define([v_maj], [2])
 m4_define([v_min], [7])
-m4_define([v_mic], [2])
+m4_define([v_mic], [3])
 ##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##
 m4_define([v_ver], [v_maj.v_min.v_mic])
 m4_define([lt_rev], m4_eval(v_maj + v_min))

From d5f485c5179c87e0e99818f96004dc146cc5e0f9 Mon Sep 17 00:00:00 2001
From: Kano <root@mutsumi.paige>
Date: Thu, 23 Aug 2012 19:20:12 +1000
Subject: [PATCH 04/17] Pools: add RollTime info to API 'stats' and 'Stats'
 button in miner.php

---
 API-README | 11 +++++++++++
 api.c      |  7 ++++++-
 miner.h    |  4 ++++
 miner.php  | 18 +++++++++++++++++-
 util.c     | 17 ++++++++++++++---
 5 files changed, 52 insertions(+), 5 deletions(-)

diff --git a/API-README b/API-README
index d4efaf6c..ff48a3c5 100644
--- a/API-README
+++ b/API-README
@@ -339,6 +339,17 @@ miner.php - an example web page to access the API
 Feature Changelog for external applications using the API:
 
 
+API V1.18
+
+Modified API commands:
+ 'stats' - add 'Work Had Roll Time', 'Work Can Roll', 'Work Had Expire',
+		'Work Roll Time' to the pool stats
+
+Modified API commands:
+ 'config' - include 'ScanTime'
+
+----------
+
 API V1.17 (cgminer v2.7.1)
 
 Added API commands:
diff --git a/api.c b/api.c
index 05d7d3ef..ca1b9429 100644
--- a/api.c
+++ b/api.c
@@ -166,7 +166,7 @@ static const char SEPARATOR = '|';
 #define SEPSTR "|"
 static const char GPUSEP = ',';
 
-static const char *APIVERSION = "1.17";
+static const char *APIVERSION = "1.18";
 static const char *DEAD = "Dead";
 static const char *SICK = "Sick";
 static const char *NOSTART = "NoStart";
@@ -1256,6 +1256,7 @@ static void minerconfig(__maybe_unused SOCKETTYPE c, __maybe_unused char *param,
 	root = api_add_const(root, "Device Code", DEVICECODE, false);
 	root = api_add_const(root, "OS", OSINFO, false);
 	root = api_add_bool(root, "Failover-Only", &opt_fail_only, false);
+	root = api_add_int(root, "ScanTime", &opt_scantime, false);
 
 	root = print_data(root, buf, isjson);
 	if (isjson)
@@ -2676,6 +2677,10 @@ static int itemstats(int i, char *id, struct cgminer_stats *stats, struct cgmine
 		root = api_add_timeval(root, "Pool Max", &(pool_stats->getwork_wait_max), false);
 		root = api_add_timeval(root, "Pool Min", &(pool_stats->getwork_wait_min), false);
 		root = api_add_double(root, "Pool Av", &(pool_stats->getwork_wait_rolling), false);
+		root = api_add_bool(root, "Work Had Roll Time", &(pool_stats->hadrolltime), false);
+		root = api_add_bool(root, "Work Can Roll", &(pool_stats->canroll), false);
+		root = api_add_bool(root, "Work Had Expire", &(pool_stats->hadexpire), false);
+		root = api_add_uint32(root, "Work Roll Time", &(pool_stats->rolltime), false);
 	}
 
 	if (extra)
diff --git a/miner.h b/miner.h
index 91fe7ce7..4f7b1d75 100644
--- a/miner.h
+++ b/miner.h
@@ -307,6 +307,10 @@ struct cgminer_pool_stats {
 	struct timeval getwork_wait_max;
 	struct timeval getwork_wait_min;
 	double getwork_wait_rolling;
+	bool hadrolltime;
+	bool canroll;
+	bool hadexpire;
+	uint32_t rolltime;
 };
 
 struct cgpu_info {
diff --git a/miner.php b/miner.php
index 198f0cf5..9c6f6949 100644
--- a/miner.php
+++ b/miner.php
@@ -84,8 +84,24 @@ $mobilesum = array(
  'DEVS+NOTIFY' => array('DEVS.MHS av', 'DEVS.Accepted', 'DEVS.Rejected', 'DEVS.Utility'),
  'POOL' => array('Accepted', 'Rejected'));
 #
+$statspage = array(
+ 'DATE' => null,
+ 'RIGS' => null,
+ 'SUMMARY' => array('Elapsed', 'MHS av', 'Found Blocks=Blks',
+			'Accepted', 'Rejected=Rej', 'Utility',
+			'Hardware Errors=HW Errs', 'Network Blocks=Net Blks',
+			'Work Utility'),
+ 'COIN' => array('*'),
+ 'STATS' => array('*'));
+#
+$statssum = array(
+ 'SUMMARY' => array('MHS av', 'Found Blocks', 'Accepted',
+			'Rejected', 'Utility', 'Hardware Errors',
+			'Work Utility'));
+#
 # customsummarypages is an array of these Custom Summary Pages
-$customsummarypages = array('Mobile' => array($mobilepage, $mobilesum));
+$customsummarypages = array('Mobile' => array($mobilepage, $mobilesum),
+ 'Stats' => array($statspage, $statssum));
 #
 $here = $_SERVER['PHP_SELF'];
 #
diff --git a/util.c b/util.c
index b104d8a3..724ffa9d 100644
--- a/util.c
+++ b/util.c
@@ -58,6 +58,9 @@ struct header_info {
 	char		*lp_path;
 	int		rolltime;
 	char		*reason;
+	bool		hadrolltime;
+	bool		canroll;
+	bool		hadexpire;
 };
 
 struct tq_ent {
@@ -157,14 +160,18 @@ static size_t resp_hdr_cb(void *ptr, size_t size, size_t nmemb, void *user_data)
 		applog(LOG_DEBUG, "HTTP hdr(%s): %s", key, val);
 
 	if (!strcasecmp("X-Roll-Ntime", key)) {
+		hi->hadrolltime = true;
 		if (!strncasecmp("N", val, 1))
 			applog(LOG_DEBUG, "X-Roll-Ntime: N found");
 		else {
+			hi->canroll = true;
+
 			/* Check to see if expire= is supported and if not, set
 			 * the rolltime to the default scantime */
-			if (strlen(val) > 7 && !strncasecmp("expire=", val, 7))
+			if (strlen(val) > 7 && !strncasecmp("expire=", val, 7)) {
 				sscanf(val + 7, "%d", &hi->rolltime);
-			else
+				hi->hadexpire = true;
+			} else
 				hi->rolltime = opt_scantime;
 			applog(LOG_DEBUG, "X-Roll-Ntime expiry set to %d", hi->rolltime);
 		}
@@ -258,7 +265,7 @@ json_t *json_rpc_call(CURL *curl, const char *url,
 {
 	long timeout = longpoll ? (60 * 60) : 60;
 	struct data_buffer all_data = {NULL, 0};
-	struct header_info hi = {NULL, 0, NULL};
+	struct header_info hi = {NULL, 0, NULL, false, false, false};
 	char len_hdr[64], user_agent_hdr[128];
 	char curl_err_str[CURL_ERROR_SIZE];
 	struct curl_slist *headers = NULL;
@@ -388,6 +395,10 @@ json_t *json_rpc_call(CURL *curl, const char *url,
 	}
 
 	*rolltime = hi.rolltime;
+	pool->cgminer_pool_stats.rolltime = hi.rolltime;
+	pool->cgminer_pool_stats.hadrolltime = hi.hadrolltime;
+	pool->cgminer_pool_stats.canroll = hi.canroll;
+	pool->cgminer_pool_stats.hadexpire = hi.hadexpire;
 
 	val = JSON_LOADS(all_data.buf, &err);
 	if (!val) {

From 4fbe5bed152a55e3f5baaa012e88c11d2f335bb8 Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Thu, 23 Aug 2012 23:25:32 +1000
Subject: [PATCH 05/17] OpenCL 1.0 does not have native atomic_add and
 extremely slow support with atom_add so detect opencl1.0 and use a non-atomic
 workaround.

---
 diablo120823.cl  | 54 +++++++++++++++++++++++++-----------------------
 diakgcn120823.cl | 54 +++++++++++++++++++++++++-----------------------
 ocl.c            |  3 +++
 phatk120823.cl   | 54 +++++++++++++++++++++++++-----------------------
 poclbm120823.cl  | 42 +++++++++++++++++++++----------------
 scrypt120823.cl  | 18 ++++++++++++++--
 6 files changed, 127 insertions(+), 98 deletions(-)

diff --git a/diablo120823.cl b/diablo120823.cl
index 4687c5bc..f4055aa6 100644
--- a/diablo120823.cl
+++ b/diablo120823.cl
@@ -1244,28 +1244,33 @@ void search(
     
 #define FOUND (0x0F)
 
+#if defined(OCL1)
+	#define SETFOUND(Xfound, Xnonce) do {	\
+		(Xfound) = output[FOUND];	\
+		output[FOUND] += 1;		\
+		output[Xfound] = Xnonce;	\
+	} while (0)
+#else
+	#define SETFOUND(Xfound, Xnonce) do {	\
+		Xfound = atomic_add(&output[FOUND], 1); \
+		output[Xfound] = Xnonce;	\
+	} while (0)
+#endif
+
 #if defined(VECTORS4)
 	bool result = any(ZA[924] == 0x136032EDU);
 
 	if (result) {
 		uint found;
 
-		if (ZA[924].x == 0x136032EDU) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = Znonce.x;
-		}
-		if (ZA[924].y == 0x136032EDU) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = Znonce.y;
-		}
-		if (ZA[924].z == 0x136032EDU) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = Znonce.z;
-		}
-		if (ZA[924].w == 0x136032EDU) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = Znonce.w;
-		}
+		if (ZA[924].x == 0x136032EDU)
+			SETFOUND(found, Znonce.x);
+		if (ZA[924].y == 0x136032EDU)
+			SETFOUND(found, Znonce.y);
+		if (ZA[924].z == 0x136032EDU)
+			SETFOUND(found, Znonce.z);
+		if (ZA[924].w == 0x136032EDU)
+			SETFOUND(found, Znonce.w);
 	}
 #elif defined(VECTORS2)
 	bool result = any(ZA[924] == 0x136032EDU);
@@ -1273,19 +1278,16 @@ void search(
 	if (result) {
 		uint found;
 
-		if (ZA[924].x == 0x136032EDU) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = Znonce.x;
-		}
-		if (ZA[924].y == 0x136032EDU) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = Znonce.y;
-		}
+		if (ZA[924].x == 0x136032EDU)
+			SETFOUND(found, Znonce.x);
+		if (ZA[924].y == 0x136032EDU)
+			SETFOUND(found, Znonce.y);
 	}
 #else
 	if (ZA[924] == 0x136032EDU) {
-		uint found = atomic_add(&output[FOUND], 1);
-		output[found] = Znonce;
+		uint found;
+
+		SETFOUND(found, Znonce);
 	}
 #endif
 }
diff --git a/diakgcn120823.cl b/diakgcn120823.cl
index d27674f6..fb4b154e 100644
--- a/diakgcn120823.cl
+++ b/diakgcn120823.cl
@@ -573,44 +573,46 @@ __kernel
 
 #define FOUND (0x0F)
 
+#if defined(OCL1)
+	#define SETFOUND(Xfound, Xnonce) do {	\
+		(Xfound) = output[FOUND];	\
+		output[FOUND] += 1;		\
+		output[Xfound] = Xnonce;	\
+	} while (0)
+#else
+	#define SETFOUND(Xfound, Xnonce) do {	\
+		Xfound = atomic_add(&output[FOUND], 1); \
+		output[Xfound] = Xnonce;	\
+	} while (0)
+#endif
+
 #ifdef VECTORS4
 	if ((V[7].x == 0x136032edU) ^ (V[7].y == 0x136032edU) ^ (V[7].z == 0x136032edU) ^ (V[7].w == 0x136032edU)) {
 		uint found;
 
-		if (V[7].x == 0x136032edU) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = nonce.x;
-		}
-		if (V[7].y == 0x136032edU) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = nonce.y;
-		}
-		if (V[7].z == 0x136032edU) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = nonce.z;
-		}
-		if (V[7].w == 0x136032edU) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = nonce.w;
-		}
+		if (V[7].x == 0x136032edU)
+			SETFOUND(found, nonce.x);
+		if (V[7].y == 0x136032edU)
+			SETFOUND(found, nonce.y);
+		if (V[7].z == 0x136032edU)
+			SETFOUND(found, nonce.z);
+		if (V[7].w == 0x136032edU)
+			SETFOUND(found, nonce.w);
 	}
 #elif defined VECTORS2
 	if ((V[7].x == 0x136032edU) + (V[7].y == 0x136032edU)) {
 		uint found;
 
-		if (V[7].x == 0x136032edU) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = nonce.x;
-		}
-		if (V[7].y == 0x136032edU) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = nonce.y;
-		}
+		if (V[7].x == 0x136032edU)
+			SETFOUND(found, nonce.x);
+		if (V[7].y == 0x136032edU)
+			SETFOUND(found, nonce.y);
 	}
 #else
 	if (V[7] == 0x136032edU) {
-		uint found = atomic_add(&output[FOUND], 1);
-		output[found] = nonce;
+		uint found;
+
+		SETFOUND(found, nonce);
 	}
 #endif
 }
diff --git a/ocl.c b/ocl.c
index 7bf606cc..fe457822 100644
--- a/ocl.c
+++ b/ocl.c
@@ -659,6 +659,9 @@ build:
 	if (clState->goffset)
 		strcat(CompilerOptions, " -D GOFFSET");
 
+	if (!clState->hasOpenCL11plus)
+		strcat(CompilerOptions, " -D OCL1");
+
 	applog(LOG_DEBUG, "CompilerOptions: %s", CompilerOptions);
 	status = clBuildProgram(clState->program, 1, &devices[gpu], CompilerOptions , NULL, NULL);
 	free(CompilerOptions);
diff --git a/phatk120823.cl b/phatk120823.cl
index cf5eb09c..4693cd8a 100644
--- a/phatk120823.cl
+++ b/phatk120823.cl
@@ -389,46 +389,48 @@ void search(	const uint state0, const uint state1, const uint state2, const uint
 
 #define FOUND (0x0F)
 
+#if defined(OCL1)
+	#define SETFOUND(Xfound, Xnonce) do {	\
+		(Xfound) = output[FOUND];	\
+		output[FOUND] += 1;		\
+		output[Xfound] = Xnonce;	\
+	} while (0)
+#else
+	#define SETFOUND(Xfound, Xnonce) do {	\
+		Xfound = atomic_add(&output[FOUND], 1); \
+		output[Xfound] = Xnonce;	\
+	} while (0)
+#endif
+
 #ifdef VECTORS4
 	bool result = W[117].x & W[117].y & W[117].z & W[117].w;
 	if (!result) {
 		uint found;
 
-		if (!W[117].x) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = W[3].x;
-		}
-		if (!W[117].y) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = W[3].y;
-		}
-		if (!W[117].z) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = W[3].z;
-		}
-		if (!W[117].w) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = W[3].w;
-		}
+		if (!W[117].x)
+			SETFOUND(found, W[3].x);
+		if (!W[117].y)
+			SETFOUND(found, W[3].y);
+		if (!W[117].z)
+			SETFOUND(found, W[3].z);
+		if (!W[117].w)
+			SETFOUND(found, W[3].w);
 	}
 #elif defined VECTORS2
 	bool result = W[117].x & W[117].y;
 	if (!result) {
 		uint found;
 
-		if (!W[117].x) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = W[3].x;
-		}
-		if (!W[117].y) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = W[3].y;
-		}
+		if (!W[117].x)
+			SETFOUND(found, W[3].x);
+		if (!W[117].y)
+			SETFOUND(found, W[3].y);
 	}
 #else
 	if (!W[117]) {
-		uint found = atomic_add(&output[FOUND], 1);
-		output[found] = W[3];
+		uint found;
+
+		SETFOUND(found, W[3]);
 	}
 #endif
 }
diff --git a/poclbm120823.cl b/poclbm120823.cl
index a02413bb..64bdb27f 100644
--- a/poclbm120823.cl
+++ b/poclbm120823.cl
@@ -1323,34 +1323,40 @@ Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
 
 #define FOUND (0x0F)
 
+#if defined(OCL1)
+	#define SETFOUND(Xfound, Xnonce) do {	\
+		(Xfound) = output[FOUND];	\
+		output[FOUND] += 1;		\
+		output[Xfound] = Xnonce;	\
+	} while (0)
+#else
+	#define SETFOUND(Xfound, Xnonce) do {	\
+		Xfound = atomic_add(&output[FOUND], 1); \
+		output[Xfound] = Xnonce;	\
+	} while (0)
+#endif
+
 #if defined(VECTORS2) || defined(VECTORS4)
 
 	if (any(Vals[2] == 0x136032edU)) {
 		uint found;
 
-		if (Vals[2].x == 0x136032edU) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = nonce.x;
-		}
-		if (Vals[2].y == 0x136032edU) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = nonce.y;
-		}
+		if (Vals[2].x == 0x136032edU)
+			SETFOUND(found, nonce.x);
+		if (Vals[2].y == 0x136032edU)
+			SETFOUND(found, nonce.y);
 #if defined(VECTORS4)
-		if (Vals[2].z == 0x136032edU) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = nonce.z;
-		}
-		if (Vals[2].w == 0x136032edU) {
-			found = atomic_add(&output[FOUND], 1);
-			output[found] = nonce.w;
-		}
+		if (Vals[2].z == 0x136032edU)
+			SETFOUND(found, nonce.z);
+		if (Vals[2].w == 0x136032edU)
+			SETFOUND(found, nonce.w);
 #endif
 	}
 #else
 	if (Vals[2] == 0x136032edU) {
-		uint found = atomic_add(&output[FOUND], 1);
-		output[found] = nonce;
+		uint found;
+
+		SETFOUND(found, nonce);
 	}
 #endif
 }
diff --git a/scrypt120823.cl b/scrypt120823.cl
index 7390d2cd..bcc0b705 100644
--- a/scrypt120823.cl
+++ b/scrypt120823.cl
@@ -684,6 +684,19 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 
 #define FOUND (0x0F)
 
+#if defined(OCL1)
+	#define SETFOUND(Xfound, Xnonce) do {	\
+		(Xfound) = output[FOUND];	\
+		output[FOUND] += 1;		\
+		output[Xfound] = Xnonce;	\
+	} while (0)
+#else
+	#define SETFOUND(Xfound, Xnonce) do {	\
+		Xfound = atomic_add(&output[FOUND], 1); \
+		output[Xfound] = Xnonce;	\
+	} while (0)
+#endif
+
 __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
 __kernel void search(__global const uint4 * restrict input,
 volatile __global uint*restrict output, __global uint4*restrict padcache,
@@ -722,8 +735,9 @@ const uint4 midstate0, const uint4 midstate16, const uint target)
 
 	bool result = (EndianSwap(ostate1.w) <= target);
 	if (result) {
-		uint found = atomic_add(&output[FOUND], 1);
-		output[found] = gid;
+		uint found;
+
+		SETFOUND(found, gid);
 	}
 }
 

From 618b3e8b116c9ac95f8bd9a24bff96369ace8bdf Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Thu, 23 Aug 2012 23:50:38 +1000
Subject: [PATCH 06/17] Track queued and staged per pool once again for future
 use.

---
 cgminer.c | 28 ++++++++++++----------------
 miner.h   |  3 +++
 2 files changed, 15 insertions(+), 16 deletions(-)

diff --git a/cgminer.c b/cgminer.c
index 29362c9b..31ab6096 100644
--- a/cgminer.c
+++ b/cgminer.c
@@ -2252,17 +2252,19 @@ static void push_curl_entry(struct curl_ent *ce, struct pool *pool)
 
 /* This is overkill, but at least we'll know accurately how much work is
  * queued to prevent ever being left without work */
-static void inc_queued(void)
+static void inc_queued(struct pool *pool)
 {
 	mutex_lock(&qd_lock);
 	total_queued++;
+	pool->queued++;
 	mutex_unlock(&qd_lock);
 }
 
-static void dec_queued(void)
+static void dec_queued(struct pool *pool)
 {
 	mutex_lock(&qd_lock);
 	total_queued--;
+	pool->queued--;
 	mutex_unlock(&qd_lock);
 }
 
@@ -2422,7 +2424,7 @@ retry:
 			lagging = true;
 		pool = ret_work->pool = select_pool(lagging);
 
-		inc_queued();
+		inc_queued(pool);
 
 		if (!ce)
 			ce = pop_curl_entry(pool);
@@ -2430,7 +2432,7 @@ retry:
 		/* Check that we haven't staged work via other threads while
 		 * waiting for a curl entry */
 		if (total_staged() >= maxq) {
-			dec_queued();
+			dec_queued(pool);
 			free_work(ret_work);
 			goto out;
 		}
@@ -2440,7 +2442,7 @@ retry:
 			/* pause, then restart work-request loop */
 			applog(LOG_DEBUG, "json_rpc_call failed on get work, retrying");
 			lagging = true;
-			dec_queued();
+			dec_queued(pool);
 			free_work(ret_work);
 			goto retry;
 		}
@@ -2725,6 +2727,7 @@ static void discard_stale(void)
 	HASH_ITER(hh, staged_work, work, tmp) {
 		if (stale_work(work, false)) {
 			HASH_DEL(staged_work, work);
+			work->pool->staged--;
 			discard_work(work);
 			stale++;
 		}
@@ -2931,9 +2934,11 @@ static bool hash_push(struct work *work)
 	pthread_cond_signal(&getq->cond);
 	mutex_unlock(stgd_lock);
 
+	work->pool->staged++;
+
 	if (work->queued) {
 		work->queued = false;
-		dec_queued();
+		dec_queued(work->pool);
 	}
 
 	return rc;
@@ -3926,17 +3931,7 @@ static void pool_resus(struct pool *pool)
 static bool queue_request(struct thr_info *thr, bool needed)
 {
 	struct workio_cmd *wc;
-	bool doq = true;
 
-	mutex_lock(&control_lock);
-	if (queued_getworks > (mining_threads + opt_queue) * 2)
-		doq = false;
-	else
-		queued_getworks++;
-	mutex_unlock(&control_lock);
-	if (!doq)
-		return true;
-	
 	/* fill out work request message */
 	wc = calloc(1, sizeof(*wc));
 	if (unlikely(!wc)) {
@@ -3981,6 +3976,7 @@ static struct work *hash_pop(const struct timespec *abstime)
 		} else
 			work = staged_work;
 		HASH_DEL(staged_work, work);
+		work->pool->staged--;
 		if (work_rollable(work))
 			staged_rollable--;
 	}
diff --git a/miner.h b/miner.h
index 91fe7ce7..2141724a 100644
--- a/miner.h
+++ b/miner.h
@@ -732,6 +732,9 @@ struct pool {
 	int solved;
 	int diff1;
 
+	int queued;
+	int staged;
+
 	bool submit_fail;
 	bool idle;
 	bool lagging;

From 37fa7d36d433f94833c6006b8447f18ec0f2b1c6 Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Fri, 24 Aug 2012 00:35:26 +1000
Subject: [PATCH 07/17] Move the decision to queue further work upstream before
 threads are spawned based on fine grained per-pool stats and increment the
 queued count immediately.

---
 cgminer.c | 67 +++++++++++++++++++++++++------------------------------
 1 file changed, 31 insertions(+), 36 deletions(-)

diff --git a/cgminer.c b/cgminer.c
index 31ab6096..856537c8 100644
--- a/cgminer.c
+++ b/cgminer.c
@@ -70,7 +70,7 @@ struct workio_cmd {
 	enum workio_commands	cmd;
 	struct thr_info		*thr;
 	struct work		*work;
-	bool			needed;
+	struct pool		*pool;
 };
 
 struct strategies strategies[] = {
@@ -187,7 +187,6 @@ int hw_errors;
 int total_accepted, total_rejected, total_diff1;
 int total_getworks, total_stale, total_discarded;
 static int total_queued, staged_rollable;
-static int queued_getworks;
 unsigned int new_blocks;
 static unsigned int work_block;
 unsigned int found_blocks;
@@ -2383,31 +2382,25 @@ out:
 	return cloned;
 }
 
+static bool queue_request(struct thr_info *thr, bool needed);
+
 static void *get_work_thread(void *userdata)
 {
 	struct workio_cmd *wc = (struct workio_cmd *)userdata;
-	int ts, tq, maxq = opt_queue + mining_threads;
 	struct pool *pool = current_pool();
 	struct work *ret_work= NULL;
 	struct curl_ent *ce = NULL;
-	bool lagging = false;
 
 	pthread_detach(pthread_self());
 
 	applog(LOG_DEBUG, "Creating extra get work thread");
 
-retry:
-	tq = global_queued();
-	ts = total_staged();
-
-	if (ts >= maxq)
-		goto out;
-
-	if (ts >= opt_queue && tq >= maxq)
-		goto out;
+	pool = wc->pool;
 
-	if (clone_available())
+	if (clone_available()) {
+		dec_queued(pool);
 		goto out;
+	}
 
 	ret_work = make_work();
 	if (wc->thr)
@@ -2419,32 +2412,19 @@ retry:
 		get_benchmark_work(ret_work);
 		ret_work->queued = true;
 	} else {
-
-		if (!ts)
-			lagging = true;
-		pool = ret_work->pool = select_pool(lagging);
-
-		inc_queued(pool);
+		ret_work->pool = wc->pool;
 
 		if (!ce)
 			ce = pop_curl_entry(pool);
 
-		/* Check that we haven't staged work via other threads while
-		 * waiting for a curl entry */
-		if (total_staged() >= maxq) {
-			dec_queued(pool);
-			free_work(ret_work);
-			goto out;
-		}
-
 		/* obtain new work from bitcoin via JSON-RPC */
 		if (!get_upstream_work(ret_work, ce->curl)) {
 			/* pause, then restart work-request loop */
 			applog(LOG_DEBUG, "json_rpc_call failed on get work, retrying");
-			lagging = true;
 			dec_queued(pool);
+			queue_request(ret_work->thr, true);
 			free_work(ret_work);
-			goto retry;
+			goto out;
 		}
 
 		ret_work->queued = true;
@@ -2463,9 +2443,6 @@ out:
 	workio_cmd_free(wc);
 	if (ce)
 		push_curl_entry(ce, pool);
-	mutex_lock(&control_lock);
-	queued_getworks--;
-	mutex_unlock(&control_lock);
 	return NULL;
 }
 
@@ -2626,8 +2603,6 @@ static struct pool *priority_pool(int choice)
 	return ret;
 }
 
-static bool queue_request(struct thr_info *thr, bool needed);
-
 void switch_pools(struct pool *selected)
 {
 	struct pool *pool, *last_pool;
@@ -3930,8 +3905,28 @@ static void pool_resus(struct pool *pool)
 
 static bool queue_request(struct thr_info *thr, bool needed)
 {
+	int ts, tq, maxq = opt_queue + mining_threads;
+	struct pool *pool, *cp;
 	struct workio_cmd *wc;
 
+	ts = total_staged();
+	tq = global_queued();
+	if (ts && ts + tq >= maxq)
+		return true;
+
+	cp = current_pool();
+	if ((!needed || opt_fail_only) && (cp->staged + cp->queued >= maxq))
+		return true;
+
+	if (needed && !ts)
+		pool = select_pool(true);
+	else
+		pool = cp;
+	if (pool->staged + pool->queued >= maxq)
+		return true;
+
+	inc_queued(pool);
+
 	/* fill out work request message */
 	wc = calloc(1, sizeof(*wc));
 	if (unlikely(!wc)) {
@@ -3941,7 +3936,7 @@ static bool queue_request(struct thr_info *thr, bool needed)
 
 	wc->cmd = WC_GET_WORK;
 	wc->thr = thr;
-	wc->needed = needed;
+	wc->pool = pool;
 
 	applog(LOG_DEBUG, "Queueing getwork request to work thread");
 

From 7992e5f3c87c649257b971e1011d0e64148b1bbc Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Fri, 24 Aug 2012 00:41:14 +1000
Subject: [PATCH 08/17] Carry the needed bool over the work command queue.

---
 cgminer.c | 4 +++-
 1 file changed, 3 insertions(+), 1 deletion(-)

diff --git a/cgminer.c b/cgminer.c
index 856537c8..ff7bd5fd 100644
--- a/cgminer.c
+++ b/cgminer.c
@@ -71,6 +71,7 @@ struct workio_cmd {
 	struct thr_info		*thr;
 	struct work		*work;
 	struct pool		*pool;
+	bool			needed;
 };
 
 struct strategies strategies[] = {
@@ -2422,7 +2423,7 @@ static void *get_work_thread(void *userdata)
 			/* pause, then restart work-request loop */
 			applog(LOG_DEBUG, "json_rpc_call failed on get work, retrying");
 			dec_queued(pool);
-			queue_request(ret_work->thr, true);
+			queue_request(ret_work->thr, wc->needed);
 			free_work(ret_work);
 			goto out;
 		}
@@ -3937,6 +3938,7 @@ static bool queue_request(struct thr_info *thr, bool needed)
 	wc->cmd = WC_GET_WORK;
 	wc->thr = thr;
 	wc->pool = pool;
+	wc->needed = needed;
 
 	applog(LOG_DEBUG, "Queueing getwork request to work thread");
 

From 9bec1e158e0e913cc9166b1f01a75e029ef9a998 Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Fri, 24 Aug 2012 01:06:17 +1000
Subject: [PATCH 09/17] Simplify macros in output kernels avoiding apparent
 loops and local variables.

---
 diablo120823.cl  | 34 ++++++++++------------------------
 diakgcn120823.cl | 34 ++++++++++------------------------
 phatk120823.cl   | 34 ++++++++++------------------------
 poclbm120823.cl  | 29 ++++++++---------------------
 scrypt120823.cl  | 18 ++++--------------
 5 files changed, 42 insertions(+), 107 deletions(-)

diff --git a/diablo120823.cl b/diablo120823.cl
index f4055aa6..b287f4fd 100644
--- a/diablo120823.cl
+++ b/diablo120823.cl
@@ -1245,49 +1245,35 @@ void search(
 #define FOUND (0x0F)
 
 #if defined(OCL1)
-	#define SETFOUND(Xfound, Xnonce) do {	\
-		(Xfound) = output[FOUND];	\
-		output[FOUND] += 1;		\
-		output[Xfound] = Xnonce;	\
-	} while (0)
+	#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
 #else
-	#define SETFOUND(Xfound, Xnonce) do {	\
-		Xfound = atomic_add(&output[FOUND], 1); \
-		output[Xfound] = Xnonce;	\
-	} while (0)
+	#define SETFOUND(Xnonce) output[atomic_add(&output[FOUND], 1)] = Xnonce
 #endif
 
 #if defined(VECTORS4)
 	bool result = any(ZA[924] == 0x136032EDU);
 
 	if (result) {
-		uint found;
-
 		if (ZA[924].x == 0x136032EDU)
-			SETFOUND(found, Znonce.x);
+			SETFOUND(Znonce.x);
 		if (ZA[924].y == 0x136032EDU)
-			SETFOUND(found, Znonce.y);
+			SETFOUND(Znonce.y);
 		if (ZA[924].z == 0x136032EDU)
-			SETFOUND(found, Znonce.z);
+			SETFOUND(Znonce.z);
 		if (ZA[924].w == 0x136032EDU)
-			SETFOUND(found, Znonce.w);
+			SETFOUND(Znonce.w);
 	}
 #elif defined(VECTORS2)
 	bool result = any(ZA[924] == 0x136032EDU);
 
 	if (result) {
-		uint found;
-
 		if (ZA[924].x == 0x136032EDU)
-			SETFOUND(found, Znonce.x);
+			SETFOUND(Znonce.x);
 		if (ZA[924].y == 0x136032EDU)
-			SETFOUND(found, Znonce.y);
+			SETFOUND(Znonce.y);
 	}
 #else
-	if (ZA[924] == 0x136032EDU) {
-		uint found;
-
-		SETFOUND(found, Znonce);
-	}
+	if (ZA[924] == 0x136032EDU)
+		SETFOUND(Znonce);
 #endif
 }
diff --git a/diakgcn120823.cl b/diakgcn120823.cl
index fb4b154e..b8e76863 100644
--- a/diakgcn120823.cl
+++ b/diakgcn120823.cl
@@ -574,45 +574,31 @@ __kernel
 #define FOUND (0x0F)
 
 #if defined(OCL1)
-	#define SETFOUND(Xfound, Xnonce) do {	\
-		(Xfound) = output[FOUND];	\
-		output[FOUND] += 1;		\
-		output[Xfound] = Xnonce;	\
-	} while (0)
+	#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
 #else
-	#define SETFOUND(Xfound, Xnonce) do {	\
-		Xfound = atomic_add(&output[FOUND], 1); \
-		output[Xfound] = Xnonce;	\
-	} while (0)
+	#define SETFOUND(Xnonce) output[atomic_add(&output[FOUND], 1)] = Xnonce
 #endif
 
 #ifdef VECTORS4
 	if ((V[7].x == 0x136032edU) ^ (V[7].y == 0x136032edU) ^ (V[7].z == 0x136032edU) ^ (V[7].w == 0x136032edU)) {
-		uint found;
-
 		if (V[7].x == 0x136032edU)
-			SETFOUND(found, nonce.x);
+			SETFOUND(nonce.x);
 		if (V[7].y == 0x136032edU)
-			SETFOUND(found, nonce.y);
+			SETFOUND(nonce.y);
 		if (V[7].z == 0x136032edU)
-			SETFOUND(found, nonce.z);
+			SETFOUND(nonce.z);
 		if (V[7].w == 0x136032edU)
-			SETFOUND(found, nonce.w);
+			SETFOUND(nonce.w);
 	}
 #elif defined VECTORS2
 	if ((V[7].x == 0x136032edU) + (V[7].y == 0x136032edU)) {
-		uint found;
-
 		if (V[7].x == 0x136032edU)
-			SETFOUND(found, nonce.x);
+			SETFOUND(nonce.x);
 		if (V[7].y == 0x136032edU)
-			SETFOUND(found, nonce.y);
+			SETFOUND(nonce.y);
 	}
 #else
-	if (V[7] == 0x136032edU) {
-		uint found;
-
-		SETFOUND(found, nonce);
-	}
+	if (V[7] == 0x136032edU)
+		SETFOUND(nonce);
 #endif
 }
diff --git a/phatk120823.cl b/phatk120823.cl
index 4693cd8a..fb1ce3c1 100644
--- a/phatk120823.cl
+++ b/phatk120823.cl
@@ -390,47 +390,33 @@ void search(	const uint state0, const uint state1, const uint state2, const uint
 #define FOUND (0x0F)
 
 #if defined(OCL1)
-	#define SETFOUND(Xfound, Xnonce) do {	\
-		(Xfound) = output[FOUND];	\
-		output[FOUND] += 1;		\
-		output[Xfound] = Xnonce;	\
-	} while (0)
+	#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
 #else
-	#define SETFOUND(Xfound, Xnonce) do {	\
-		Xfound = atomic_add(&output[FOUND], 1); \
-		output[Xfound] = Xnonce;	\
-	} while (0)
+	#define SETFOUND(Xnonce) output[atomic_add(&output[FOUND], 1)] = Xnonce
 #endif
 
 #ifdef VECTORS4
 	bool result = W[117].x & W[117].y & W[117].z & W[117].w;
 	if (!result) {
-		uint found;
-
 		if (!W[117].x)
-			SETFOUND(found, W[3].x);
+			SETFOUND(W[3].x);
 		if (!W[117].y)
-			SETFOUND(found, W[3].y);
+			SETFOUND(W[3].y);
 		if (!W[117].z)
-			SETFOUND(found, W[3].z);
+			SETFOUND(W[3].z);
 		if (!W[117].w)
-			SETFOUND(found, W[3].w);
+			SETFOUND(W[3].w);
 	}
 #elif defined VECTORS2
 	bool result = W[117].x & W[117].y;
 	if (!result) {
-		uint found;
-
 		if (!W[117].x)
-			SETFOUND(found, W[3].x);
+			SETFOUND(W[3].x);
 		if (!W[117].y)
-			SETFOUND(found, W[3].y);
+			SETFOUND(W[3].y);
 	}
 #else
-	if (!W[117]) {
-		uint found;
-
-		SETFOUND(found, W[3]);
-	}
+	if (!W[117])
+		SETFOUND(W[3]);
 #endif
 }
diff --git a/poclbm120823.cl b/poclbm120823.cl
index 64bdb27f..9ae2ee94 100644
--- a/poclbm120823.cl
+++ b/poclbm120823.cl
@@ -1324,39 +1324,26 @@ Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
 #define FOUND (0x0F)
 
 #if defined(OCL1)
-	#define SETFOUND(Xfound, Xnonce) do {	\
-		(Xfound) = output[FOUND];	\
-		output[FOUND] += 1;		\
-		output[Xfound] = Xnonce;	\
-	} while (0)
+	#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
 #else
-	#define SETFOUND(Xfound, Xnonce) do {	\
-		Xfound = atomic_add(&output[FOUND], 1); \
-		output[Xfound] = Xnonce;	\
-	} while (0)
+	#define SETFOUND(Xnonce) output[atomic_add(&output[FOUND], 1)] = Xnonce
 #endif
 
 #if defined(VECTORS2) || defined(VECTORS4)
-
 	if (any(Vals[2] == 0x136032edU)) {
-		uint found;
-
 		if (Vals[2].x == 0x136032edU)
-			SETFOUND(found, nonce.x);
+			SETFOUND(nonce.x);
 		if (Vals[2].y == 0x136032edU)
-			SETFOUND(found, nonce.y);
+			SETFOUND(nonce.y);
 #if defined(VECTORS4)
 		if (Vals[2].z == 0x136032edU)
-			SETFOUND(found, nonce.z);
+			SETFOUND(nonce.z);
 		if (Vals[2].w == 0x136032edU)
-			SETFOUND(found, nonce.w);
+			SETFOUND(nonce.w);
 #endif
 	}
 #else
-	if (Vals[2] == 0x136032edU) {
-		uint found;
-
-		SETFOUND(found, nonce);
-	}
+	if (Vals[2] == 0x136032edU)
+		SETFOUND(nonce);
 #endif
 }
diff --git a/scrypt120823.cl b/scrypt120823.cl
index bcc0b705..4b884583 100644
--- a/scrypt120823.cl
+++ b/scrypt120823.cl
@@ -685,16 +685,9 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 #define FOUND (0x0F)
 
 #if defined(OCL1)
-	#define SETFOUND(Xfound, Xnonce) do {	\
-		(Xfound) = output[FOUND];	\
-		output[FOUND] += 1;		\
-		output[Xfound] = Xnonce;	\
-	} while (0)
+	#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
 #else
-	#define SETFOUND(Xfound, Xnonce) do {	\
-		Xfound = atomic_add(&output[FOUND], 1); \
-		output[Xfound] = Xnonce;	\
-	} while (0)
+	#define SETFOUND(Xnonce) output[atomic_add(&output[FOUND], 1)] = Xnonce
 #endif
 
 __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
@@ -734,11 +727,8 @@ const uint4 midstate0, const uint4 midstate16, const uint target)
 	SHA256(&ostate0,&ostate1, tmp0, tmp1, (uint4)(0x80000000U, 0U, 0U, 0U), (uint4)(0U, 0U, 0U, 0x300U));
 
 	bool result = (EndianSwap(ostate1.w) <= target);
-	if (result) {
-		uint found;
-
-		SETFOUND(found, gid);
-	}
+	if (result)
+		SETFOUND(gid);
 }
 
 /*-

From 0e0093e6020ed955a177af8eb576e42de5078129 Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Fri, 24 Aug 2012 01:23:54 +1000
Subject: [PATCH 10/17] Select pool regardless of whether we're lagging or not,
 and don't queue another request in switch pool to avoid infinite recursion.

---
 cgminer.c | 7 +------
 1 file changed, 1 insertion(+), 6 deletions(-)

diff --git a/cgminer.c b/cgminer.c
index ff7bd5fd..d7396b1e 100644
--- a/cgminer.c
+++ b/cgminer.c
@@ -2678,8 +2678,6 @@ void switch_pools(struct pool *selected)
 	mutex_lock(&lp_lock);
 	pthread_cond_broadcast(&lp_cond);
 	mutex_unlock(&lp_lock);
-
-	queue_request(NULL, false);
 }
 
 static void discard_work(struct work *work)
@@ -3919,10 +3917,7 @@ static bool queue_request(struct thr_info *thr, bool needed)
 	if ((!needed || opt_fail_only) && (cp->staged + cp->queued >= maxq))
 		return true;
 
-	if (needed && !ts)
-		pool = select_pool(true);
-	else
-		pool = cp;
+	pool = select_pool(needed && !ts);
 	if (pool->staged + pool->queued >= maxq)
 		return true;
 

From cf6033cb0a0ed59114c3f34f91a1b1447da1f6a3 Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Fri, 24 Aug 2012 01:26:55 +1000
Subject: [PATCH 11/17] Update NEWS.

---
 NEWS | 14 ++++++++++++++
 1 file changed, 14 insertions(+)

diff --git a/NEWS b/NEWS
index adf0a2e3..06f328dd 100644
--- a/NEWS
+++ b/NEWS
@@ -1,3 +1,17 @@
+Version 2.7.4 - August 23, 2012
+
+- Perform select_pool even when not lagging to allow it to switch back if needed
+to the primary.
+- Simplify macros in output kernels avoiding apparent loops and local variables.
+- Carry the needed bool over the work command queue.
+- Move the decision to queue further work upstream before threads are spawned
+based on fine grained per-pool stats and increment the queued count immediately.
+- Track queued and staged per pool once again for future use.
+- OpenCL 1.0 does not have native atomic_add and extremely slow support with
+atom_add so detect opencl1.0 and use a non-atomic workaround.
+- Pools: add RollTime info to API 'stats' and 'Stats' button in miner.php
+
+
 Version 2.7.3 - August 22, 2012
 
 - Minimise the number of getwork threads we generate.

From 7fe8142cbf4aaf5b81db9987dd5225ae723489b9 Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Fri, 24 Aug 2012 01:27:16 +1000
Subject: [PATCH 12/17] Bump version to 2.7.4.

---
 configure.ac | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/configure.ac b/configure.ac
index 7ee4885d..4197da9c 100644
--- a/configure.ac
+++ b/configure.ac
@@ -2,7 +2,7 @@
 ##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##
 m4_define([v_maj], [2])
 m4_define([v_min], [7])
-m4_define([v_mic], [3])
+m4_define([v_mic], [4])
 ##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##
 m4_define([v_ver], [v_maj.v_min.v_mic])
 m4_define([lt_rev], m4_eval(v_maj + v_min))

From 1b7db5bc9c774d18e76ef416799c39fe43d867fe Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Mon, 27 Aug 2012 09:47:55 +1000
Subject: [PATCH 13/17] thr is always NULL going into queue_request now.

---
 cgminer.c | 18 +++++++-----------
 1 file changed, 7 insertions(+), 11 deletions(-)

diff --git a/cgminer.c b/cgminer.c
index d7396b1e..004989b4 100644
--- a/cgminer.c
+++ b/cgminer.c
@@ -2383,7 +2383,7 @@ out:
 	return cloned;
 }
 
-static bool queue_request(struct thr_info *thr, bool needed);
+static bool queue_request(bool needed);
 
 static void *get_work_thread(void *userdata)
 {
@@ -2404,10 +2404,7 @@ static void *get_work_thread(void *userdata)
 	}
 
 	ret_work = make_work();
-	if (wc->thr)
-		ret_work->thr = wc->thr;
-	else
-		ret_work->thr = NULL;
+	ret_work->thr = NULL;
 
 	if (opt_benchmark) {
 		get_benchmark_work(ret_work);
@@ -2423,7 +2420,7 @@ static void *get_work_thread(void *userdata)
 			/* pause, then restart work-request loop */
 			applog(LOG_DEBUG, "json_rpc_call failed on get work, retrying");
 			dec_queued(pool);
-			queue_request(ret_work->thr, wc->needed);
+			queue_request(wc->needed);
 			free_work(ret_work);
 			goto out;
 		}
@@ -2711,7 +2708,7 @@ static void discard_stale(void)
 	if (stale) {
 		applog(LOG_DEBUG, "Discarded %d stales that didn't match current hash", stale);
 		while (stale-- > 0)
-			queue_request(NULL, false);
+			queue_request(false);
 	}
 }
 
@@ -3902,7 +3899,7 @@ static void pool_resus(struct pool *pool)
 		switch_pools(NULL);
 }
 
-static bool queue_request(struct thr_info *thr, bool needed)
+static bool queue_request(bool needed)
 {
 	int ts, tq, maxq = opt_queue + mining_threads;
 	struct pool *pool, *cp;
@@ -3931,7 +3928,6 @@ static bool queue_request(struct thr_info *thr, bool needed)
 	}
 
 	wc->cmd = WC_GET_WORK;
-	wc->thr = thr;
 	wc->pool = pool;
 	wc->needed = needed;
 
@@ -3974,7 +3970,7 @@ static struct work *hash_pop(const struct timespec *abstime)
 	}
 	mutex_unlock(stgd_lock);
 
-	queue_request(NULL, false);
+	queue_request(false);
 
 	return work;
 }
@@ -5724,7 +5720,7 @@ begin_bench:
 #endif
 
 	for (i = 0; i < mining_threads + opt_queue; i++)
-		queue_request(NULL, false);
+		queue_request(false);
 
 	/* main loop - simply wait for workio thread to exit. This is not the
 	 * normal exit path and only occurs should the workio_thread die

From d1683f75c9d033509a3b2cf80317eb147c203f9d Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Mon, 27 Aug 2012 09:55:19 +1000
Subject: [PATCH 14/17] needed flag is currently always false in queue_request.
 Remove it for now.

---
 cgminer.c | 18 ++++++++----------
 1 file changed, 8 insertions(+), 10 deletions(-)

diff --git a/cgminer.c b/cgminer.c
index 004989b4..e56e592c 100644
--- a/cgminer.c
+++ b/cgminer.c
@@ -71,7 +71,6 @@ struct workio_cmd {
 	struct thr_info		*thr;
 	struct work		*work;
 	struct pool		*pool;
-	bool			needed;
 };
 
 struct strategies strategies[] = {
@@ -2383,7 +2382,7 @@ out:
 	return cloned;
 }
 
-static bool queue_request(bool needed);
+static bool queue_request(void);
 
 static void *get_work_thread(void *userdata)
 {
@@ -2420,7 +2419,7 @@ static void *get_work_thread(void *userdata)
 			/* pause, then restart work-request loop */
 			applog(LOG_DEBUG, "json_rpc_call failed on get work, retrying");
 			dec_queued(pool);
-			queue_request(wc->needed);
+			queue_request();
 			free_work(ret_work);
 			goto out;
 		}
@@ -2708,7 +2707,7 @@ static void discard_stale(void)
 	if (stale) {
 		applog(LOG_DEBUG, "Discarded %d stales that didn't match current hash", stale);
 		while (stale-- > 0)
-			queue_request(false);
+			queue_request();
 	}
 }
 
@@ -3899,7 +3898,7 @@ static void pool_resus(struct pool *pool)
 		switch_pools(NULL);
 }
 
-static bool queue_request(bool needed)
+static bool queue_request(void)
 {
 	int ts, tq, maxq = opt_queue + mining_threads;
 	struct pool *pool, *cp;
@@ -3911,10 +3910,10 @@ static bool queue_request(bool needed)
 		return true;
 
 	cp = current_pool();
-	if ((!needed || opt_fail_only) && (cp->staged + cp->queued >= maxq))
+	if (cp->staged + cp->queued >= maxq)
 		return true;
 
-	pool = select_pool(needed && !ts);
+	pool = select_pool(false);
 	if (pool->staged + pool->queued >= maxq)
 		return true;
 
@@ -3929,7 +3928,6 @@ static bool queue_request(bool needed)
 
 	wc->cmd = WC_GET_WORK;
 	wc->pool = pool;
-	wc->needed = needed;
 
 	applog(LOG_DEBUG, "Queueing getwork request to work thread");
 
@@ -3970,7 +3968,7 @@ static struct work *hash_pop(const struct timespec *abstime)
 	}
 	mutex_unlock(stgd_lock);
 
-	queue_request(false);
+	queue_request();
 
 	return work;
 }
@@ -5720,7 +5718,7 @@ begin_bench:
 #endif
 
 	for (i = 0; i < mining_threads + opt_queue; i++)
-		queue_request(false);
+		queue_request();
 
 	/* main loop - simply wait for workio thread to exit. This is not the
 	 * normal exit path and only occurs should the workio_thread die

From 4a210d4eff61b710d95d0e4a7f95f89ae9d04858 Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Mon, 27 Aug 2012 10:02:53 +1000
Subject: [PATCH 15/17] Only clear the pool lagging flag if we're staging work
 faster than we're using it.

---
 cgminer.c | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/cgminer.c b/cgminer.c
index e56e592c..0f882726 100644
--- a/cgminer.c
+++ b/cgminer.c
@@ -4085,7 +4085,10 @@ retry:
 	pool = work_heap->pool;
 	/* If we make it here we have succeeded in getting fresh work */
 	if (!work_heap->mined) {
-		pool_tclear(pool, &pool->lagging);
+		/* Only clear the lagging flag if we are staging them at a
+		 * rate faster then we're using them */
+		if (pool->lagging && total_staged())
+			pool_tclear(pool, &pool->lagging);
 		if (pool_tclear(pool, &pool->idle))
 			pool_resus(pool);
 	}

From 579c1299c681abfafa2eea2ca69866c4f2f321b1 Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Mon, 27 Aug 2012 10:10:50 +1000
Subject: [PATCH 16/17] There is no need to try to switch pools in select_pool
 since the current pool is actually not affected by the choice of pool to get
 work from.

---
 cgminer.c | 8 +++-----
 1 file changed, 3 insertions(+), 5 deletions(-)

diff --git a/cgminer.c b/cgminer.c
index 0f882726..8dc1deb7 100644
--- a/cgminer.c
+++ b/cgminer.c
@@ -1955,11 +1955,9 @@ static inline struct pool *select_pool(bool lagging)
 	if (pool_strategy == POOL_BALANCE)
 		return select_balanced(cp);
 
-	if (pool_strategy != POOL_LOADBALANCE && (!lagging || opt_fail_only)) {
-		if (cp->prio != 0)
-			switch_pools(NULL);
-		pool = current_pool();
-	} else
+	if (pool_strategy != POOL_LOADBALANCE && (!lagging || opt_fail_only))
+		pool = cp;
+	else
 		pool = NULL;
 
 	while (!pool) {

From b7687588181744191ec3d7ff690c588a85c20fbf Mon Sep 17 00:00:00 2001
From: Con Kolivas <kernel@kolivas.org>
Date: Mon, 27 Aug 2012 10:15:48 +1000
Subject: [PATCH 17/17] Test for lagging once more in queue_request to enable
 work to leak to backup pools.

---
 cgminer.c | 6 ++++--
 1 file changed, 4 insertions(+), 2 deletions(-)

diff --git a/cgminer.c b/cgminer.c
index 8dc1deb7..f677a60b 100644
--- a/cgminer.c
+++ b/cgminer.c
@@ -3901,6 +3901,7 @@ static bool queue_request(void)
 	int ts, tq, maxq = opt_queue + mining_threads;
 	struct pool *pool, *cp;
 	struct workio_cmd *wc;
+	bool lagging;
 
 	ts = total_staged();
 	tq = global_queued();
@@ -3908,10 +3909,11 @@ static bool queue_request(void)
 		return true;
 
 	cp = current_pool();
-	if (cp->staged + cp->queued >= maxq)
+	lagging = !opt_fail_only && cp->lagging && !ts && cp->queued >= maxq;
+	if (!lagging && cp->staged + cp->queued >= maxq)
 		return true;
 
-	pool = select_pool(false);
+	pool = select_pool(lagging);
 	if (pool->staged + pool->queued >= maxq)
 		return true;