Merge branch 'master' into bfgminer

configure.ac

@@ -333,7 +333,7 @@ AC_DEFINE_UNQUOTED([CGMINER_PREFIX], ["$prefix/bin"], [Path to bfgminer install]
 
 AC_DEFINE_UNQUOTED([PHATK_KERNNAME], ["phatk120223"], [Filename for phatk kernel])
 AC_DEFINE_UNQUOTED([POCLBM_KERNNAME], ["poclbm120327"], [Filename for poclbm kernel])
-AC_DEFINE_UNQUOTED([DIAKGCN_KERNNAME], ["diakgcn120223"], [Filename for diakgcn kernel])
+AC_DEFINE_UNQUOTED([DIAKGCN_KERNNAME], ["diakgcn120427"], [Filename for diakgcn kernel])
 AC_DEFINE_UNQUOTED([DIABLO_KERNNAME], ["diablo120328"], [Filename for diablo kernel])
 
 

diakgcn120223.cl → diakgcn120427.cl

@@ -1,4 +1,4 @@
-// DiaKGCN 16-03-2012 - OpenCL kernel by Diapolo
+// DiaKGCN 27-04-2012 - OpenCL kernel by Diapolo
 //
 // Parts and / or ideas for this kernel are based upon the public-domain poclbm project, the phatk kernel by Phateus and the DiabloMiner kernel by DiabloD3.
 // The kernel was rewritten by me (Diapolo) and is still public-domain!
@@ -33,7 +33,9 @@
 __kernel
 	__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
 	void search(	
+#ifndef GOFFSET
 			const u base,
+#endif
 			const uint PreVal0, const uint PreVal4,
 			const uint H1, const uint D1A, const uint B1, const uint C1,
 			const uint F1, const uint G1, const uint C1addK5, const uint B1addK6, const uint PreVal0addK7,
@@ -56,7 +58,11 @@ __kernel
 #elif defined VECTORS2
 	const u nonce = (uint)(get_local_id(0)) * 2U + (uint)(get_group_id(0)) * (uint)(WORKVEC) + base;
 #else
-	const u nonce = (uint)(get_local_id(0)) + (uint)(get_group_id(0)) * (uint)(WORKSIZE) + base;
+	#ifdef GOFFSET
+		const u nonce = (uint)(get_global_id(0));
+	#else
+		const u nonce = (uint)(get_local_id(0)) + (uint)(get_group_id(0)) * (uint)(WORKSIZE) + base;
+	#endif
 #endif
 
 	V[0] = PreVal0 + nonce;

driver-opencl.c

@@ -1096,15 +1096,17 @@ static cl_int queue_diakgcn_kernel(_clState *clState, dev_blk_ctx *blk,
 				   __maybe_unused cl_uint threads)
 {
 	cl_kernel *kernel = &clState->kernel;
-	cl_uint vwidth = clState->vwidth;
-	unsigned int i, num = 0;
+	unsigned int num = 0;
 	cl_int status = 0;
-	uint *nonces;
 
-	nonces = alloca(sizeof(uint) * vwidth);
-	for (i = 0; i < vwidth; i++)
-		nonces[i] = blk->nonce + i;
-	CL_SET_VARG(vwidth, nonces);
+	if (!clState->goffset) {
+		cl_uint vwidth = clState->vwidth;
+		uint *nonces = alloca(sizeof(uint) * vwidth);
+		unsigned int i;
+		for (i = 0; i < vwidth; i++)
+			nonces[i] = blk->nonce + i;
+		CL_SET_VARG(vwidth, nonces);
+	}
 
 	CL_SET_BLKARG(PreVal0);
 	CL_SET_BLKARG(PreVal4_2);

miner.c

@@ -1657,8 +1657,7 @@ static bool submit_upstream_work(const struct work *work)
 		if (isblock)
 			found_blocks++;
 		hash32 = (uint32_t *)(work->hash);
-		sprintf(hashshow, "%08lx.%08lx.%08lx%s",
-			(unsigned long)(hash32[7]), (unsigned long)(hash32[6]), (unsigned long)(hash32[5]),
+		sprintf(hashshow, "%08lx.%08lx%s", (unsigned long)(hash32[6]), (unsigned long)(hash32[5]),
 			isblock ? " BLOCK!" : "");
 #endif
 	}
@@ -1675,11 +1674,11 @@ static bool submit_upstream_work(const struct work *work)
 		applog(LOG_DEBUG, "PROOF OF WORK RESULT: true (yay!!!)");
 		if (!QUIET) {
 			if (total_pools > 1)
-				applog(LOG_NOTICE, "Accepted %s %s %d thread %d pool %d",
-				       hashshow, cgpu->api->name, cgpu->device_id, thr_id, work->pool->pool_no);
+				applog(LOG_NOTICE, "Accepted %s %s %d pool %d",
+				       hashshow, cgpu->api->name, cgpu->device_id, work->pool->pool_no);
 			else
-				applog(LOG_NOTICE, "Accepted %s %s %d thread %d",
-				       hashshow, cgpu->api->name, cgpu->device_id, thr_id);
+				applog(LOG_NOTICE, "Accepted %s %s %d",
+				       hashshow, cgpu->api->name, cgpu->device_id);
 		}
 		sharelog("accept", work);
 		if (opt_shares && total_accepted >= opt_shares) {
@@ -3821,10 +3820,6 @@ static void convert_to_work(json_t *val, bool rolltime, struct pool *pool)
 	struct work *work, *work_clone;
 	bool rc;
 
-	/* Don't use as work if we have failover-only enabled */
-	if (pool != cp && opt_fail_only)
-		return;
-
 	work = make_work();
 
 	rc = work_decode(json_object_get(val, "result"), work);
@@ -3846,6 +3841,12 @@ static void convert_to_work(json_t *val, bool rolltime, struct pool *pool)
 	 * allows testwork to know whether LP discovered the block or not. */
 	test_work_current(work);
 
+	/* Don't use as work if we have failover-only enabled */
+	if (pool != cp && opt_fail_only) {
+		free_work(work);
+		return;
+	}
+
 	work_clone = make_work();
 	memcpy(work_clone, work, sizeof(struct work));
 	while (reuse_work(work)) {
@@ -3912,6 +3913,9 @@ retry_pool:
 		}
 	}
 
+	/* Any longpoll from any pool is enough for this to be true */
+	have_longpoll = true;
+
 	if (cp == pool)
 		applog(LOG_WARNING, "Long-polling activated for %s", pool->lp_url);
 	else

ocl.c

@@ -598,7 +598,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
 		gpus[gpu].vwidth = preferred_vwidth;
 	}
 
-	if ((clState->chosen_kernel == KL_POCLBM || clState->chosen_kernel == KL_DIABLO) &&
+	if ((clState->chosen_kernel == KL_POCLBM || clState->chosen_kernel == KL_DIABLO || clState->chosen_kernel == KL_DIAKGCN) &&
 		clState->vwidth == 1 && clState->hasOpenCL11plus)
 			clState->goffset = true;