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/NEWS b/NEWS
index dbd8b9cd..06f328dd 100644
--- a/NEWS
+++ b/NEWS
@@ -1,3 +1,22 @@
+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.
+
+
 Version 2.7.2 - August 22, 2012
 
 - Pick worksize 256 with Cypress if none is specified.
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/cgminer.c b/cgminer.c
index 5b15e859..f677a60b 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[] = {
@@ -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) {
@@ -2251,17 +2249,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);
 }
 
@@ -2380,68 +2380,46 @@ out:
 	return cloned;
 }
 
+static bool queue_request(void);
+
 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;
+	pool = wc->pool;
 
-	if (ts >= opt_queue && tq >= maxq)
-		goto out;
-
-	if (clone_available())
+	if (clone_available()) {
+		dec_queued(pool);
 		goto out;
+	}
 
 	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);
 		ret_work->queued = true;
 	} else {
-
-		if (!ts)
-			lagging = true;
-		pool = ret_work->pool = select_pool(lagging);
-
-		inc_queued();
+		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();
-			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();
+			dec_queued(pool);
+			queue_request();
 			free_work(ret_work);
-			goto retry;
+			goto out;
 		}
 
 		ret_work->queued = true;
@@ -2620,8 +2598,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;
@@ -2696,8 +2672,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)
@@ -2721,6 +2695,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++;
 		}
@@ -2730,7 +2705,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();
 	}
 }
 
@@ -2927,9 +2902,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;
@@ -3919,9 +3896,28 @@ static void pool_resus(struct pool *pool)
 		switch_pools(NULL);
 }
 
-static bool queue_request(struct thr_info *thr, bool needed)
+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();
+	if (ts && ts + tq >= maxq)
+		return true;
+
+	cp = current_pool();
+	lagging = !opt_fail_only && cp->lagging && !ts && cp->queued >= maxq;
+	if (!lagging && cp->staged + cp->queued >= maxq)
+		return true;
+
+	pool = select_pool(lagging);
+	if (pool->staged + pool->queued >= maxq)
+		return true;
+
+	inc_queued(pool);
 
 	/* fill out work request message */
 	wc = calloc(1, sizeof(*wc));
@@ -3931,8 +3927,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");
 
@@ -3967,12 +3962,13 @@ 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--;
 	}
 	mutex_unlock(stgd_lock);
 
-	queue_request(NULL, false);
+	queue_request();
 
 	return work;
 }
@@ -4089,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);
 	}
@@ -5722,7 +5721,7 @@ begin_bench:
 #endif
 
 	for (i = 0; i < mining_threads + opt_queue; i++)
-		queue_request(NULL, 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
diff --git a/configure.ac b/configure.ac
index ae5d677f..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], [2])
+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))
diff --git a/diablo120823.cl b/diablo120823.cl
index 4687c5bc..b287f4fd 100644
--- a/diablo120823.cl
+++ b/diablo120823.cl
@@ -1244,48 +1244,36 @@ void search(
     
 #define FOUND (0x0F)
 
+#if defined(OCL1)
+	#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
+#else
+	#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) {
-			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(Znonce.x);
+		if (ZA[924].y == 0x136032EDU)
+			SETFOUND(Znonce.y);
+		if (ZA[924].z == 0x136032EDU)
+			SETFOUND(Znonce.z);
+		if (ZA[924].w == 0x136032EDU)
+			SETFOUND(Znonce.w);
 	}
 #elif defined(VECTORS2)
 	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].x == 0x136032EDU)
+			SETFOUND(Znonce.x);
+		if (ZA[924].y == 0x136032EDU)
+			SETFOUND(Znonce.y);
 	}
 #else
-	if (ZA[924] == 0x136032EDU) {
-		uint found = atomic_add(&output[FOUND], 1);
-		output[found] = Znonce;
-	}
+	if (ZA[924] == 0x136032EDU)
+		SETFOUND(Znonce);
 #endif
 }
diff --git a/diakgcn120823.cl b/diakgcn120823.cl
index d27674f6..b8e76863 100644
--- a/diakgcn120823.cl
+++ b/diakgcn120823.cl
@@ -573,44 +573,32 @@ __kernel
 
 #define FOUND (0x0F)
 
+#if defined(OCL1)
+	#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
+#else
+	#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) {
-			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(nonce.x);
+		if (V[7].y == 0x136032edU)
+			SETFOUND(nonce.y);
+		if (V[7].z == 0x136032edU)
+			SETFOUND(nonce.z);
+		if (V[7].w == 0x136032edU)
+			SETFOUND(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(nonce.x);
+		if (V[7].y == 0x136032edU)
+			SETFOUND(nonce.y);
 	}
 #else
-	if (V[7] == 0x136032edU) {
-		uint found = atomic_add(&output[FOUND], 1);
-		output[found] = nonce;
-	}
+	if (V[7] == 0x136032edU)
+		SETFOUND(nonce);
 #endif
 }
diff --git a/miner.h b/miner.h
index 91fe7ce7..13f4960c 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 {
@@ -732,6 +736,9 @@ struct pool {
 	int solved;
 	int diff1;
 
+	int queued;
+	int staged;
+
 	bool submit_fail;
 	bool idle;
 	bool lagging;
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/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..fb1ce3c1 100644
--- a/phatk120823.cl
+++ b/phatk120823.cl
@@ -389,46 +389,34 @@ void search(	const uint state0, const uint state1, const uint state2, const uint
 
 #define FOUND (0x0F)
 
+#if defined(OCL1)
+	#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
+#else
+	#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) {
-			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(W[3].x);
+		if (!W[117].y)
+			SETFOUND(W[3].y);
+		if (!W[117].z)
+			SETFOUND(W[3].z);
+		if (!W[117].w)
+			SETFOUND(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(W[3].x);
+		if (!W[117].y)
+			SETFOUND(W[3].y);
 	}
 #else
-	if (!W[117]) {
-		uint found = atomic_add(&output[FOUND], 1);
-		output[found] = W[3];
-	}
+	if (!W[117])
+		SETFOUND(W[3]);
 #endif
 }
diff --git a/poclbm120823.cl b/poclbm120823.cl
index a02413bb..9ae2ee94 100644
--- a/poclbm120823.cl
+++ b/poclbm120823.cl
@@ -1323,34 +1323,27 @@ Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
 
 #define FOUND (0x0F)
 
-#if defined(VECTORS2) || defined(VECTORS4)
+#if defined(OCL1)
+	#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
+#else
+	#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) {
-			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(nonce.x);
+		if (Vals[2].y == 0x136032edU)
+			SETFOUND(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(nonce.z);
+		if (Vals[2].w == 0x136032edU)
+			SETFOUND(nonce.w);
 #endif
 	}
 #else
-	if (Vals[2] == 0x136032edU) {
-		uint found = atomic_add(&output[FOUND], 1);
-		output[found] = nonce;
-	}
+	if (Vals[2] == 0x136032edU)
+		SETFOUND(nonce);
 #endif
 }
diff --git a/scrypt120823.cl b/scrypt120823.cl
index 7390d2cd..4b884583 100644
--- a/scrypt120823.cl
+++ b/scrypt120823.cl
@@ -684,6 +684,12 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 
 #define FOUND (0x0F)
 
+#if defined(OCL1)
+	#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
+#else
+	#define SETFOUND(Xnonce) output[atomic_add(&output[FOUND], 1)] = Xnonce
+#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,
@@ -721,10 +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 = atomic_add(&output[FOUND], 1);
-		output[found] = gid;
-	}
+	if (result)
+		SETFOUND(gid);
 }
 
 /*-
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) {