Merge branch 'cg_merges_20130524b' into bfgminer

Makefile.am

@@ -6,6 +6,7 @@ EXTRA_DIST	= example.conf m4/gnulib-cache.m4 linux-usb-bfgminer \
 		  API.class API.java api-example.c windows-build.txt \
 		  bitstreams/* README.FPGA README.RPC README.scrypt \
 		  README.GPU \
+		  hexdump.c \
                   api-example.py
 
 SUBDIRS		= lib ccan
@@ -19,7 +20,7 @@ bin_SCRIPTS	= $(top_srcdir)/*.cl
 bfgminer_LDFLAGS	= $(PTHREAD_FLAGS)
 bfgminer_LDADD	= $(DLOPEN_FLAGS) @LIBCURL_LIBS@ @JANSSON_LIBS@ @PTHREAD_LIBS@ \
 		  @NCURSES_LIBS@ @PDCURSES_LIBS@ @WS2_LIBS@ \
-		  @UDEV_LIBS@ @LIBUSB_LIBS@ \
+		  @UDEV_LIBS@ @LIBUSB_LIBS@ @MM_LIBS@ \
 		  @MATH_LIBS@ lib/libgnu.a ccan/libccan.a
 bfgminer_CPPFLAGS = -I$(top_builddir)/lib -I$(top_srcdir)/lib @LIBUSB_CFLAGS@ @LIBCURL_CFLAGS@
 
@@ -118,7 +119,7 @@ bfgminer_SOURCES += driver-cairnsmore.c
 endif
 
 if HAS_AVALON
-bfgminer_SOURCES += driver-avalon.c
+bfgminer_SOURCES += driver-avalon.c driver-avalon.h
 endif
 
 if HAS_MODMINER

README

@@ -641,6 +641,14 @@ A: Start BFGMiner with your regular commands and add -D -T --verbose and provide
 the full startup output and a summary of your hardware, operating system, ATI
 driver version and ATI stream version.
 
+Q: Can I mine with BFGMiner on a Mac?
+A: BFGMiner will compile on OS X, but the performance of GPU mining is
+compromised due to the OpenCL implementation on OS X, there is no temperature or
+fanspeed monitoring, and the cooling design of most Macs, despite having
+powerful GPUs, will usually not cope with constant usage leading to a high risk
+of thermal damage. It is highly recommended not to mine on a Mac unless it is to
+a USB device.
+
 Q: My network gets slower and slower and then dies for a minute?
 A; Try the --net-delay option.
 

README.GPU

@@ -357,7 +357,8 @@ Q: Which AMD driver is the best?
 A: Unfortunately AMD has a history of having quite a few releases with issues
 when it comes to mining, either in terms of breaking mining, increasing CPU
 usage or very low hashrates. Only experimentation can tell you for sure, but
-some good releases were 11.6, 11.12, 12.4 and 12.8
+some good releases were 11.6, 11.12, 12.4 and 12.8. Note that older cards may
+not work with the newer drivers.
 
 Q: I have multiple SDKs installed, can I choose which one it uses?
 A: Run bfgminer with the -n option and it will list all the platforms currently
@@ -407,6 +408,10 @@ Q: Can I mine with Nvidia or Intel GPUs?
 A: Yes, but their hashrate is very poor and likely you'll be using much more
 energy than you'll be earning in coins.
 
+Q: Can I mine on both Nvidia and AMD GPUs at the same time?
+A: No, you must run one instance of BFGMiner with the --gpu-platform option for
+each.
+
 Q: Can I mine on Linux without running Xorg?
 A: With Nvidia you can, but with AMD you cannot.
 
@@ -422,12 +427,45 @@ Q: My scrypt hashrate is high but the pool reports only a tiny proportion of my
 hashrate?
 A: You are generating garbage hashes due to your choice of settings. Try
 decreasing your intensity, do not increase the number of gpu-threads, and
-consider adding system RAM to match your GPU ram.
+consider adding system RAM to match your GPU ram. You may also be using a bad
+combination of driver and/or SDK.
 
 Q: Scrypt fails to initialise the kernel every time?
 A: Your parameters are too high. Don't add GPU threads, don't set intensity too
 high, decrease thread concurrency. See the README.scrypt for a lot more help.
 
+Q: BFGMiner stops mining (or my GPUs go DEAD) and I can't close it?
+A: Once the driver has crashed, there is no way for BFGMiner to close cleanly.
+You will have to kill it, and depending on how corrupted your driver state has
+gotten, you may even need to reboot. Windows is known to reset drivers when they
+fail and BFGMiner will be stuck trying to use the old driver instance.
+
+Q: I can't get any monitoring of temperatures or fanspeed with BFGMiner when I
+start it remotely?
+A: With Linux, make sure to export the DISPLAY variable. On Windows, you cannot
+access these monitoring values via RDP. This should work with TightVNC or
+TeamViewer, though.
+
+Q: I change my GPU engine/memory/voltage and BFGMiner reports back no change?
+A: BFGMiner asks the GPU using the ATI Display Library to change settings, but
+the driver and hardware are free to do what it wants with that query, including
+ignoring it. Some GPUs are locked with one or more of those properties as well.
+
+Q: I have multiple GPUs and although many devices show up, it appears to be
+working only on one GPU splitting it up.
+A: Your driver setup is failing to properly use the accessory GPUs. Your driver
+may be configured wrong or you have a driver version that needs a dummy plug on
+all the GPUs that aren't connected to a monitor.
+
+Q: I have some random GPU performance related problem not addressed above.
+A: Seriously, it's the driver and/or SDK. Uninstall them and start again,
+noting there is no clean way to uninstall them so you have to use extra tools
+or do it manually.
+
+Q: Do I need to recompile after updating my driver/SDK?
+A: No. The software is unchanged regardless of which driver/SDK/ADL version you
+are running.
+
 Q: Should I use crossfire/SLI?
 A: It does not benefit mining at all and depending on the GPU may actually
 worsen performance.

README.scrypt

@@ -218,8 +218,8 @@ without crashing the GPU, you will have to use a lower memclock.
 Then, and only then, bother trying to increase intensity further.
 
 My final settings were:
---gpu-engine 1157  --gpu-memclock 1900 -I 20
-for a hashrate of 725kH.
+--gpu-engine 1141  --gpu-memclock 1875 -I 20
+for a hashrate of 745kH.
 
 Note I did not bother setting a thread concurrency. Once you have the magic
 endpoint, look at what tc was chosen by the bin file generated and then hard

api-example.c

@@ -138,7 +138,7 @@
 static const char SEPARATOR = '|';
 static const char COMMA = ',';
 static const char EQ = '=';
-static int ONLY = 0;
+static int ONLY;
 
 void display(char *buf)
 {

api.c

@@ -161,9 +161,6 @@ static const char *SCRYPTSTR = "scrypt";
 static const char *SHA256STR = "sha256";
 
 static const char *DEVICECODE = ""
-#ifdef USE_AVALON
-			"AVA "
-#endif
 #ifdef HAVE_OPENCL
 			"GPU "
 #endif
@@ -173,6 +170,9 @@ static const char *DEVICECODE = ""
 #ifdef USE_ICARUS
 			"ICA "
 #endif
+#ifdef USE_AVALON
+			"AVA "
+#endif
 #ifdef USE_X6500
 			"XBS "
 #endif
@@ -836,7 +836,7 @@ static struct api_data *api_add_data_full(struct api_data *root, char *name, enu
 
 	api_data = (struct api_data *)malloc(sizeof(struct api_data));
 
-	api_data->name = name;
+	api_data->name = strdup(name);
 	api_data->type = type;
 
 	if (root == NULL) {
@@ -1143,6 +1143,7 @@ static struct api_data *print_data(struct api_data *root, char *buf, bool isjson
 
 		buf = strchr(buf, '\0');
 
+		free(root->name);
 		if (root->type == API_JSON)
 			json_decref((json_t *)root->data);
 		if (root->data_was_malloc)
@@ -3277,7 +3278,7 @@ static void send_result(struct io_data *io_data, SOCKETTYPE c, bool isjson)
 		n = send(c, buf, tosend, 0);
 
 		if (SOCKETFAIL(n)) {
-			if (errno == EAGAIN || errno == EWOULDBLOCK)
+			if (sock_blocks())
 				continue;
 
 			applog(LOG_WARNING, "API: send (%d) failed: %s", tosend, SOCKERRMSG);

autogen.sh

@@ -14,5 +14,5 @@ echo 'Running autoreconf -if...'
 	cd "${bs_dir}"
 	rm -rf autom4te.cache
 	rm -f aclocal.m4 ltmain.sh
-	autoreconf -if
+	autoreconf -if ${AC_FLAGS}
 )

ccan/Makefile.am

@@ -1,3 +1,4 @@
 noinst_LIBRARIES	= libccan.a
 
 libccan_a_SOURCES	= compiler/compiler.h opt/helpers.c opt/opt.c opt/opt.h opt/parse.c opt/private.h opt/usage.c typesafe_cb/typesafe_cb.h
+libccan_a_CPPFLAGS	= -I$(top_srcdir)

compat.h

@@ -79,12 +79,14 @@ struct tm *localtime_convert(time_t t)
 #endif
 
 #ifndef HAVE_NANOSLEEP
+extern void cgtime(struct timeval *);
+
 static inline int nanosleep(const struct timespec *req, struct timespec *rem)
 {
 	struct timeval tstart;
 	DWORD msecs;
 
-	gettimeofday(&tstart, NULL);
+	cgtime(&tstart);
 	msecs = (req->tv_sec * 1000) + ((999999 + req->tv_nsec) / 1000000);
 
 	if (SleepEx(msecs, true) == WAIT_IO_COMPLETION) {
@@ -97,7 +99,7 @@ static inline int nanosleep(const struct timespec *req, struct timespec *rem)
 				++tdone.tv_sec;
 			}
 
-			gettimeofday(&tnow, NULL);
+			cgtime(&tnow);
 			if (timercmp(&tnow, &tdone, >))
 				return 0;
 			timersub(&tdone, &tnow, &tleft);

configure.ac

@@ -66,6 +66,7 @@ have_cygwin=false
 have_win32=false
 DLOPEN_FLAGS="-ldl"
 WS2_LIBS=""
+MM_LIBS=""
 MATH_LIBS="-lm"
 
 case $target in
@@ -88,6 +89,7 @@ case $target in
     have_win32=true
     DLOPEN_FLAGS=""
     WS2_LIBS="-lws2_32"
+    MM_LIBS="-lwinmm"
     AC_DEFINE([_WIN32_WINNT], [0x0501], "WinNT version for XP+ support")
     AC_DEFINE([FD_SETSIZE], [4096], [Maximum sockets before fd_set overflows])
     ;;
@@ -742,7 +744,7 @@ AC_DEFINE_UNQUOTED([PHATK_KERNNAME], ["phatk121016"], [Filename for phatk kernel
 AC_DEFINE_UNQUOTED([POCLBM_KERNNAME], ["poclbm130302"], [Filename for poclbm kernel])
 AC_DEFINE_UNQUOTED([DIAKGCN_KERNNAME], ["diakgcn121016"], [Filename for diakgcn kernel])
 AC_DEFINE_UNQUOTED([DIABLO_KERNNAME], ["diablo130302"], [Filename for diablo kernel])
-AC_DEFINE_UNQUOTED([SCRYPT_KERNNAME], ["scrypt130302"], [Filename for scrypt kernel])
+AC_DEFINE_UNQUOTED([SCRYPT_KERNNAME], ["scrypt130511"], [Filename for scrypt kernel])
 
 
 AC_SUBST(PTHREAD_FLAGS)
@@ -752,6 +754,7 @@ AC_SUBST(NCURSES_CPPFLAGS)
 AC_SUBST(NCURSES_LIBS)
 AC_SUBST(PDCURSES_LIBS)
 AC_SUBST(WS2_LIBS)
+AC_SUBST(MM_LIBS)
 AC_SUBST(MATH_LIBS)
 AC_SUBST(UDEV_LIBS)
 AC_SUBST(YASM_FMT)

deviceapi.c

@@ -580,7 +580,7 @@ out:
 	return NULL;
 }
 
-bool add_cgpu(struct cgpu_info*cgpu)
+bool add_cgpu(struct cgpu_info *cgpu)
 {
 	int lpcount;
 	
@@ -632,6 +632,10 @@ bool add_cgpu(struct cgpu_info*cgpu)
 	}
 	
 	wr_unlock(&devices_lock);
+
+	mutex_lock(&stats_lock);
+	cgpu->last_device_valid_work = time(NULL);
+	mutex_unlock(&stats_lock);
 	
 	return true;
 }

driver-avalon.c

@@ -17,10 +17,10 @@
 #include <stdio.h>
 #include <sys/time.h>
 #include <sys/types.h>
-#include <sys/select.h>
 #include <dirent.h>
 #include <unistd.h>
 #ifndef WIN32
+  #include <sys/select.h>
   #include <termios.h>
   #include <sys/stat.h>
   #include <fcntl.h>
@@ -28,6 +28,7 @@
     #define O_CLOEXEC 0
   #endif
 #else
+  #include "compat.h"
   #include <windows.h>
   #include <io.h>
 #endif
@@ -38,9 +39,10 @@
 #include "fpgautils.h"
 #include "driver-avalon.h"
 #include "hexdump.c"
+#include "util.h"
 
 static int option_offset = -1;
-struct avalon_info **avalon_info;
+struct avalon_info **avalon_infos;
 struct device_drv avalon_drv;
 
 static int avalon_init_task(struct avalon_task *at,
@@ -172,7 +174,7 @@ static int avalon_send_task(int fd, const struct avalon_task *at,
 	buf[4] = tt;
 #endif
 	if (likely(avalon)) {
-		info = avalon_info[avalon->device_id];
+		info = avalon_infos[avalon->device_id];
 		delay = nr_len * 10 * 1000000000ULL;
 		delay = delay / info->baud;
 	}
@@ -180,7 +182,7 @@ static int avalon_send_task(int fd, const struct avalon_task *at,
 	if (at->reset)
 		nr_len = 1;
 	if (opt_debug) {
-		applog(LOG_DEBUG, "Avalon: Sent(%d):", nr_len);
+		applog(LOG_DEBUG, "Avalon: Sent(%u):", (unsigned int)nr_len);
 		hexdump((uint8_t *)buf, nr_len);
 	}
 	ret = write(fd, buf, nr_len);
@@ -202,30 +204,40 @@ static int avalon_send_task(int fd, const struct avalon_task *at,
 	return AVA_SEND_BUFFER_EMPTY;
 }
 
-static int avalon_gets(int fd, uint8_t *buf, int read_count,
-		       struct thr_info *thr, struct timeval *tv_finish)
+static inline int avalon_gets(int fd, uint8_t *buf, struct thr_info *thr,
+		       struct timeval *tv_finish)
 {
-	ssize_t ret = 0;
-	int rc = 0;
 	int read_amount = AVALON_READ_SIZE;
 	bool first = true;
+	ssize_t ret = 0;
 
 	while (true) {
-		struct timeval timeout = {0, 100000};
+		struct timeval timeout;
 		fd_set rd;
 
+		if (unlikely(thr->work_restart)) {
+			applog(LOG_DEBUG, "Avalon: Work restart");
+			return AVA_GETS_RESTART;
+		}
+
+		timeout.tv_sec = 0;
+		timeout.tv_usec = 100000;
+
 		FD_ZERO(&rd);
 		FD_SET(fd, &rd);
 		ret = select(fd + 1, &rd, NULL, NULL, &timeout);
-		if (unlikely(ret < 0))
+		if (unlikely(ret < 0)) {
+			applog(LOG_ERR, "Avalon: Error %d on select in avalon_gets", errno);
 			return AVA_GETS_ERROR;
+		}
 		if (ret) {
 			ret = read(fd, buf, read_amount);
-			if (unlikely(ret < 0))
+			if (unlikely(ret < 0)) {
+				applog(LOG_ERR, "Avalon: Error %d on read in avalon_gets", errno);
 				return AVA_GETS_ERROR;
+			}
 			if (likely(first)) {
-				if (likely(tv_finish))
-					gettimeofday(tv_finish, NULL);
+				cgtime(tv_finish);
 				first = false;
 			}
 			if (likely(ret >= read_amount))
@@ -235,43 +247,23 @@ static int avalon_gets(int fd, uint8_t *buf, int read_count,
 			continue;
 		}
 
-		rc++;
-		if (rc >= read_count) {
-			if (opt_debug) {
-				applog(LOG_WARNING,
-				       "Avalon: No data in %.2f seconds",
-				       (float)rc/(float)AVALON_TIME_FACTOR);
-			}
-			return AVA_GETS_TIMEOUT;
-		}
-
-		if (thr && thr->work_restart) {
-			if (opt_debug) {
-				applog(LOG_WARNING,
-				       "Avalon: Work restart at %.2f seconds",
-				       (float)(rc)/(float)AVALON_TIME_FACTOR);
-			}
+		if (unlikely(thr->work_restart)) {
+			applog(LOG_DEBUG, "Avalon: Work restart");
 			return AVA_GETS_RESTART;
 		}
+
+		return AVA_GETS_TIMEOUT;
 	}
 }
 
 static int avalon_get_result(int fd, struct avalon_result *ar,
 			     struct thr_info *thr, struct timeval *tv_finish)
 {
-	struct cgpu_info *avalon;
-	struct avalon_info *info;
 	uint8_t result[AVALON_READ_SIZE];
-	int ret, read_count = AVALON_RESET_FAULT_DECISECONDS * AVALON_TIME_FACTOR;
-
-	if (likely(thr)) {
-		avalon = thr->cgpu;
-		info = avalon_info[avalon->device_id];
-		read_count = info->read_count;
-	}
+	int ret;
 
 	memset(result, 0, AVALON_READ_SIZE);
-	ret = avalon_gets(fd, result, read_count, thr, tv_finish);
+	ret = avalon_gets(fd, result, thr, tv_finish);
 
 	if (ret == AVA_GETS_OK) {
 		if (opt_debug) {
@@ -284,35 +276,65 @@ static int avalon_get_result(int fd, struct avalon_result *ar,
 	return ret;
 }
 
-static int avalon_decode_nonce(struct thr_info *thr, struct work **work,
-			       struct avalon_result *ar, uint32_t *nonce)
+static bool avalon_decode_nonce(struct thr_info *thr, struct avalon_result *ar,
+				uint32_t *nonce)
 {
 	struct cgpu_info *avalon;
 	struct avalon_info *info;
-	int avalon_get_work_count, i;
-
-	if (unlikely(!work))
-		return -1;
+	struct work *work;
 
 	avalon = thr->cgpu;
-	info = avalon_info[avalon->device_id];
-	avalon_get_work_count = info->miner_count;
+	if (unlikely(!avalon->works))
+		return false;
 
-	for (i = 0; i < avalon_get_work_count; i++) {
-		if (work[i] &&
-		    !memcmp(ar->data, work[i]->data + 64, 12) &&
-		    !memcmp(ar->midstate, work[i]->midstate, 32))
-			break;
-	}
-	if (i == avalon_get_work_count)
-		return -1;
+	work = find_queued_work_bymidstate(avalon, (char *)ar->midstate, 32,
+					   (char *)ar->data, 64, 12);
+	if (!work)
+		return false;
 
-	info->matching_work[i]++;
+	info = avalon_infos[avalon->device_id];
+	info->matching_work[work->subid]++;
 	*nonce = htole32(ar->nonce);
+	submit_nonce(thr, work, *nonce);
 
-	applog(LOG_DEBUG, "Avalon: match to work[%d](%p): %d",i, work[i],
-	       info->matching_work[i]);
-	return i;
+	return true;
+}
+
+static void avalon_get_reset(int fd, struct avalon_result *ar)
+{
+	int read_amount = AVALON_READ_SIZE;
+	uint8_t result[AVALON_READ_SIZE];
+	struct timeval timeout = {1, 0};
+	ssize_t ret = 0, offset = 0;
+	fd_set rd;
+
+	memset(result, 0, AVALON_READ_SIZE);
+	memset(ar, 0, AVALON_READ_SIZE);
+	FD_ZERO(&rd);
+	FD_SET(fd, &rd);
+	ret = select(fd + 1, &rd, NULL, NULL, &timeout);
+	if (unlikely(ret < 0)) {
+		applog(LOG_WARNING, "Avalon: Error %d on select in avalon_get_reset", errno);
+		return;
+	}
+	if (!ret) {
+		applog(LOG_WARNING, "Avalon: Timeout on select in avalon_get_reset");
+		return;
+	}
+	do {
+		ret = read(fd, result + offset, read_amount);
+		if (unlikely(ret < 0)) {
+			applog(LOG_WARNING, "Avalon: Error %d on read in avalon_get_reset", errno);
+			return;
+		}
+		read_amount -= ret;
+		offset += ret;
+	} while (read_amount > 0);
+	if (opt_debug) {
+		applog(LOG_DEBUG, "Avalon: get:");
+		hexdump((uint8_t *)result, AVALON_READ_SIZE);
+	}
+	memcpy((uint8_t *)ar, result, AVALON_READ_SIZE);
 }
 
 static int avalon_reset(int fd, struct avalon_result *ar)
@@ -333,7 +355,7 @@ static int avalon_reset(int fd, struct avalon_result *ar)
 	if (ret == AVA_SEND_ERROR)
 		return 1;
 
-	avalon_get_result(fd, ar, NULL, NULL);
+	avalon_get_reset(fd, ar);
 
 	buf = (uint8_t *)ar;
 	/* Sometimes there is one extra 0 byte for some reason in the buffer,
@@ -367,7 +389,7 @@ static void avalon_idle(struct cgpu_info *avalon)
 	struct avalon_task at;
 
 	int fd = avalon->device_fd;
-	struct avalon_info *info = avalon_info[avalon->device_id];
+	struct avalon_info *info = avalon_infos[avalon->device_id];
 	int avalon_get_work_count = info->miner_count;
 
 	i = 0;
@@ -560,19 +582,19 @@ static bool avalon_detect_one(const char *devpath)
 		   * return false; */
 	}
 	
-	avalon_info = realloc(avalon_info,
-			      sizeof(struct avalon_info *) *
-			      (total_devices + 1));
+	avalon_infos = realloc(avalon_infos,
+			       sizeof(struct avalon_info *) *
+			       (total_devices + 1));
 
 	applog(LOG_INFO, "Avalon Detect: Found at %s, mark as %d",
 	       devpath, avalon->device_id);
 
-	avalon_info[avalon->device_id] = (struct avalon_info *)
+	avalon_infos[avalon->device_id] = (struct avalon_info *)
 		malloc(sizeof(struct avalon_info));
-	if (unlikely(!(avalon_info[avalon->device_id])))
-		quit(1, "Failed to malloc avalon_info");
+	if (unlikely(!(avalon_infos[avalon->device_id])))
+		quit(1, "Failed to malloc avalon_infos");
 
-	info = avalon_info[avalon->device_id];
+	info = avalon_infos[avalon->device_id];
 
 	memset(info, 0, sizeof(struct avalon_info));
 
@@ -580,8 +602,6 @@ static bool avalon_detect_one(const char *devpath)
 	info->miner_count = miner_count;
 	info->asic_count = asic_count;
 	info->timeout = timeout;
-	info->read_count = ((float)info->timeout * AVALON_HASH_TIME_FACTOR *
-			    AVALON_TIME_FACTOR) / (float)info->miner_count;
 
 	info->fan_pwm = AVALON_DEFAULT_FAN_MIN_PWM;
 	info->temp_max = 0;
@@ -595,14 +615,11 @@ static bool avalon_detect_one(const char *devpath)
 	info->temp_old = 0;
 	info->frequency = frequency;
 
-	/* Do something for failed reset ? */
-	if (0) {
-		/* Set asic to idle mode after detect */
-		avalon_idle(avalon);
-		avalon->device_fd = -1;
+	/* Set asic to idle mode after detect */
+	avalon_idle(avalon);
+	avalon->device_fd = -1;
 
-		avalon_close(fd);
-	}
+	avalon_close(fd);
 	return true;
 }
 
@@ -623,7 +640,7 @@ static void avalon_init(struct cgpu_info *avalon)
 
 	avalon->device_fd = -1;
 	fd = avalon_open(avalon->device_path,
-			     avalon_info[avalon->device_id]->baud);
+			     avalon_infos[avalon->device_id]->baud);
 	if (unlikely(fd == -1)) {
 		applog(LOG_ERR, "Avalon: Failed to open on %s",
 		       avalon->device_path);
@@ -636,7 +653,6 @@ static void avalon_init(struct cgpu_info *avalon)
 		return;
 	}
 
-	avalon->status = LIFE_INIT;
 	avalon->device_fd = fd;
 	__avalon_init(avalon);
 }
@@ -644,43 +660,53 @@ static void avalon_init(struct cgpu_info *avalon)
 static bool avalon_prepare(struct thr_info *thr)
 {
 	struct cgpu_info *avalon = thr->cgpu;
+	struct avalon_info *info = avalon_infos[avalon->device_id];
 	struct timeval now;
 
+	free(avalon->works);
+	avalon->works = calloc(info->miner_count * sizeof(struct work *),
+			       AVALON_ARRAY_SIZE);
+	if (!avalon->works)
+		quit(1, "Failed to calloc avalon works in avalon_prepare");
 	if (avalon->device_fd == -1)
 		avalon_init(avalon);
 	else
 		__avalon_init(avalon);
 
-	gettimeofday(&now, NULL);
+	cgtime(&now);
 	get_datestamp(avalon->init, &now);
 	return true;
 }
 
-static void avalon_free_work(struct thr_info *thr, struct work **work)
+static void avalon_free_work(struct thr_info *thr)
 {
 	struct cgpu_info *avalon;
 	struct avalon_info *info;
+	struct work **works;
 	int i;
 
-	if (unlikely(!work))
-		return;
-
 	avalon = thr->cgpu;
-	info = avalon_info[avalon->device_id];
+	avalon->queued = 0;
+	if (unlikely(!avalon->works))
+		return;
+	works = avalon->works;
+	info = avalon_infos[avalon->device_id];
 
-	for (i = 0; i < info->miner_count; i++)
-		if (likely(work[i])) {
-			free_work(work[i]);
-			work[i] = NULL;
+	for (i = 0; i < info->miner_count * 4; i++) {
+		if (works[i]) {
+			work_completed(avalon, works[i]);
+			works[i] = NULL;
 		}
+	}
 }
 
 static void do_avalon_close(struct thr_info *thr)
 {
 	struct avalon_result ar;
 	struct cgpu_info *avalon = thr->cgpu;
-	struct avalon_info *info = avalon_info[avalon->device_id];
+	struct avalon_info *info = avalon_infos[avalon->device_id];
 
+	avalon_free_work(thr);
 	nmsleep(1000);
 	avalon_reset(avalon->device_fd, &ar);
 	avalon_idle(avalon);
@@ -688,16 +714,10 @@ static void do_avalon_close(struct thr_info *thr)
 	avalon->device_fd = -1;
 
 	info->no_matching_work = 0;
-	avalon_free_work(thr, info->bulk0);
-	avalon_free_work(thr, info->bulk1);
-	avalon_free_work(thr, info->bulk2);
-	avalon_free_work(thr, info->bulk3);
 }
 
 static inline void record_temp_fan(struct avalon_info *info, struct avalon_result *ar, float *temp_avg)
 {
-	int max;
-
 	info->fan0 = ar->fan0 * AVALON_FAN_FACTOR;
 	info->fan1 = ar->fan1 * AVALON_FAN_FACTOR;
 	info->fan2 = ar->fan2 * AVALON_FAN_FACTOR;
@@ -718,24 +738,14 @@ static inline void record_temp_fan(struct avalon_info *info, struct avalon_resul
 		info->temp2 = 0 - ((~ar->temp2 & 0x7f) + 1);
 	}
 
-	*temp_avg = info->temp2;
-
-	max = info->temp_max;
-	if (info->temp0 > max)
-		max = info->temp0;
-	if (info->temp1 > max)
-		max = info->temp1;
-	if (info->temp2 > max)
-		max = info->temp2;
-	if (max >= 100) {	/* FIXME: fix the root cause on fpga controller firmware */
-		if (opt_debug) {
-			applog(LOG_DEBUG, "Avalon: temp_max: %d", max);
-			hexdump((uint8_t *)ar, AVALON_READ_SIZE);
-		}
-		return;
-	}
+	*temp_avg = info->temp2 > info->temp1 ? info->temp2 : info->temp1;
 
-	info->temp_max = max;
+	if (info->temp0 > info->temp_max)
+		info->temp_max = info->temp0;
+	if (info->temp1 > info->temp_max)
+		info->temp_max = info->temp1;
+	if (info->temp2 > info->temp_max)
+		info->temp_max = info->temp2;
 }
 
 static inline void adjust_fan(struct avalon_info *info)
@@ -756,29 +766,61 @@ static inline void adjust_fan(struct avalon_info *info)
 	}
 }
 
-static int64_t avalon_scanhash(struct thr_info *thr, struct work **work,
-			       __maybe_unused int64_t max_nonce)
+/* We use a replacement algorithm to only remove references to work done from
+ * the buffer when we need the extra space for new work. */
+static bool avalon_fill(struct cgpu_info *avalon)
+{
+	int subid, slot, mc = avalon_infos[avalon->device_id]->miner_count;
+	struct work *work;
+
+	if (avalon->queued >= mc)
+		return true;
+	work = get_queued(avalon);
+	if (unlikely(!work))
+		return false;
+	subid = avalon->queued++;
+	work->subid = subid;
+	slot = avalon->work_array * mc + subid;
+	if (likely(avalon->works[slot]))
+		work_completed(avalon, avalon->works[slot]);
+	avalon->works[slot] = work;
+	if (avalon->queued >= mc)
+		return true;
+	return false;
+}
+
+static void avalon_rotate_array(struct cgpu_info *avalon)
+{
+	avalon->queued = 0;
+	if (++avalon->work_array >= AVALON_ARRAY_SIZE)
+		avalon->work_array = 0;
+}
+
+static int64_t avalon_scanhash(struct thr_info *thr)
 {
 	struct cgpu_info *avalon;
-	int fd, ret, full;
+	struct work **works;
+	int fd, ret = AVA_GETS_OK, full;
 
 	struct avalon_info *info;
 	struct avalon_task at;
 	struct avalon_result ar;
-	int i, work_i0, work_i1, work_i2, work_i3;
+	int i;
 	int avalon_get_work_count;
+	int start_count, end_count;
 
 	struct timeval tv_start, tv_finish, elapsed;
 	uint32_t nonce;
 	int64_t hash_count;
 	static int first_try = 0;
-	int result_count, result_wrong;
+	int result_wrong;
 
 	avalon = thr->cgpu;
-	info = avalon_info[avalon->device_id];
+	works = avalon->works;
+	info = avalon_infos[avalon->device_id];
 	avalon_get_work_count = info->miner_count;
 
-	if (unlikely(avalon->device_fd == -1))
+	if (unlikely(avalon->device_fd == -1)) {
 		if (!avalon_prepare(thr)) {
 			applog(LOG_ERR, "AVA%i: Comms error(open)",
 			       avalon->device_id);
@@ -786,35 +828,25 @@ static int64_t avalon_scanhash(struct thr_info *thr, struct work **work,
 			/* fail the device if the reopen attempt fails */
 			return -1;
 		}
+	}
 	fd = avalon->device_fd;
 #ifndef WIN32
 	tcflush(fd, TCOFLUSH);
 #endif
 
-	for (i = 0; i < avalon_get_work_count; i++) {
-		info->bulk0[i] = info->bulk1[i];
-		info->bulk1[i] = info->bulk2[i];
-		info->bulk2[i] = info->bulk3[i];
-		info->bulk3[i] = work[i];
-		applog(LOG_DEBUG, "Avalon: bulk0/1/2 buffer [%d]: %p, %p, %p, %p",
-		       i, info->bulk0[i], info->bulk1[i], info->bulk2[i], info->bulk3[i]);
-	}
-
-	i = 0;
+	start_count = avalon->work_array * avalon_get_work_count;
+	end_count = start_count + avalon_get_work_count;
+	i = start_count;
 	while (true) {
 		avalon_init_task(&at, 0, 0, info->fan_pwm,
 				 info->timeout, info->asic_count,
 				 info->miner_count, 1, 0, info->frequency);
-		avalon_create_task(&at, work[i]);
+		avalon_create_task(&at, works[i]);
 		ret = avalon_send_task(fd, &at, avalon);
 		if (unlikely(ret == AVA_SEND_ERROR ||
 			     (ret == AVA_SEND_BUFFER_EMPTY &&
-			      (i + 1 == avalon_get_work_count) &&
+			      (i + 1 == end_count) &&
 			      first_try))) {
-			avalon_free_work(thr, info->bulk0);
-			avalon_free_work(thr, info->bulk1);
-			avalon_free_work(thr, info->bulk2);
-			avalon_free_work(thr, info->bulk3);
 			do_avalon_close(thr);
 			applog(LOG_ERR, "AVA%i: Comms error(buffer)",
 			       avalon->device_id);
@@ -824,12 +856,13 @@ static int64_t avalon_scanhash(struct thr_info *thr, struct work **work,
 			avalon_init(avalon);
 			return 0;	/* This should never happen */
 		}
-		if (ret == AVA_SEND_BUFFER_EMPTY && (i + 1 == avalon_get_work_count)) {
+		if (ret == AVA_SEND_BUFFER_EMPTY && (i + 1 == end_count)) {
 			first_try = 1;
+			avalon_rotate_array(avalon);
 			return 0xffffffff;
 		}
 
-		work[i]->blk.nonce = 0xffffffff;
+		works[i]->blk.nonce = 0xffffffff;
 
 		if (ret == AVA_SEND_BUFFER_FULL)
 			break;
@@ -840,14 +873,11 @@ static int64_t avalon_scanhash(struct thr_info *thr, struct work **work,
 		first_try = 0;
 
 	elapsed.tv_sec = elapsed.tv_usec = 0;
-	gettimeofday(&tv_start, NULL);
+	cgtime(&tv_start);
 
-	result_count = 0;
 	result_wrong = 0;
 	hash_count = 0;
 	while (true) {
-		work_i0 = work_i1 = work_i2 = work_i3 = -1;
-
 		full = avalon_buffer_full(fd);
 		applog(LOG_DEBUG, "Avalon: Buffer full: %s",
 		       ((full == AVA_BUFFER_FULL) ? "Yes" : "No"));
@@ -856,194 +886,96 @@ static int64_t avalon_scanhash(struct thr_info *thr, struct work **work,
 
 		ret = avalon_get_result(fd, &ar, thr, &tv_finish);
 		if (unlikely(ret == AVA_GETS_ERROR)) {
-			avalon_free_work(thr, info->bulk0);
-			avalon_free_work(thr, info->bulk1);
-			avalon_free_work(thr, info->bulk2);
-			avalon_free_work(thr, info->bulk3);
 			do_avalon_close(thr);
 			applog(LOG_ERR,
 			       "AVA%i: Comms error(read)", avalon->device_id);
 			dev_error(avalon, REASON_DEV_COMMS_ERROR);
 			return 0;
 		}
+		if (unlikely(ret == AVA_GETS_RESTART))
+			break;
 		if (unlikely(ret == AVA_GETS_TIMEOUT)) {
 			timersub(&tv_finish, &tv_start, &elapsed);
 			applog(LOG_DEBUG, "Avalon: no nonce in (%ld.%06lds)",
-			       elapsed.tv_sec, elapsed.tv_usec);
+			       (long)elapsed.tv_sec, (long)elapsed.tv_usec);
 			continue;
 		}
-		if (unlikely(ret == AVA_GETS_RESTART)) {
-			avalon_free_work(thr, info->bulk0);
-			avalon_free_work(thr, info->bulk1);
-			avalon_free_work(thr, info->bulk2);
-			avalon_free_work(thr, info->bulk3);
-			break;
+
+		if (!avalon_decode_nonce(thr, &ar, &nonce)) {
+			info->no_matching_work++;
+			result_wrong++;
+
+			if (unlikely(result_wrong >= avalon_get_work_count))
+				break;
+
+			if (opt_debug) {
+				timersub(&tv_finish, &tv_start, &elapsed);
+				applog(LOG_DEBUG,"Avalon: no matching work: %d"
+				" (%ld.%06lds)", info->no_matching_work,
+				(long)elapsed.tv_sec, (long)elapsed.tv_usec);
+			}
+			continue;
 		}
-		result_count++;
-
-		work_i0 = avalon_decode_nonce(thr, info->bulk0, &ar, &nonce);
-		if (work_i0 < 0) {
-			work_i1 = avalon_decode_nonce(thr, info->bulk1, &ar, &nonce);
-			if (work_i1 < 0) {
-				work_i2 = avalon_decode_nonce(thr, info->bulk2, &ar, &nonce);
-				if (work_i2 < 0) {
-					work_i3 = avalon_decode_nonce(thr, info->bulk3, &ar, &nonce);
-					if (work_i3 < 0) {
-						info->no_matching_work++;
-						result_wrong++;
-
-						if (opt_debug) {
-							timersub(&tv_finish, &tv_start, &elapsed);
-							applog(LOG_DEBUG,"Avalon: no matching work: %d"
-							" (%ld.%06lds)", info->no_matching_work,
-							elapsed.tv_sec, elapsed.tv_usec);
-						}
-						continue;
-					} else
-						submit_nonce(thr, info->bulk3[work_i3], nonce);
-				} else
-					submit_nonce(thr, info->bulk2[work_i2], nonce);
-			} else
-				submit_nonce(thr, info->bulk1[work_i1], nonce);
-		} else
-			submit_nonce(thr, info->bulk0[work_i0], nonce);
-
-		hash_count += nonce;
+
+		hash_count += 0xffffffff;
 		if (opt_debug) {
 			timersub(&tv_finish, &tv_start, &elapsed);
 			applog(LOG_DEBUG,
-			       "Avalon: nonce = 0x%08x = 0x%08llx hashes "
-			       "(%ld.%06lds)", nonce, hash_count,
+			       "Avalon: nonce = 0x%08x = 0x%08"PRIx64" hashes "
+			       "(%ld.%06lds)", nonce, (uint64_t)hash_count,
 			       elapsed.tv_sec, elapsed.tv_usec);
 		}
 	}
-	if (result_wrong && result_count == result_wrong) {
-		/* This mean FPGA controller give all wrong result
-		 * try to reset the Avalon */
-		avalon_free_work(thr, info->bulk0);
-		avalon_free_work(thr, info->bulk1);
-		avalon_free_work(thr, info->bulk2);
-		avalon_free_work(thr, info->bulk3);
+	if (hash_count && avalon->results < AVALON_ARRAY_SIZE)
+		avalon->results++;
+	if (unlikely((result_wrong >= avalon_get_work_count) ||
+	    (!hash_count && ret != AVA_GETS_RESTART && --avalon->results < 0))) {
+		/* Look for all invalid results, or consecutive failure
+		 * to generate any results suggesting the FPGA
+		 * controller has screwed up. */
 		do_avalon_close(thr);
 		applog(LOG_ERR,
-		       "AVA%i: FPGA controller mess up", avalon->device_id);
+			"AVA%i: FPGA controller messed up, %d wrong results",
+			avalon->device_id, result_wrong);
 		dev_error(avalon, REASON_DEV_COMMS_ERROR);
-		do_avalon_close(thr);
 		nmsleep(1000);
 		avalon_init(avalon);
 		return 0;
 	}
 
-	avalon_free_work(thr, info->bulk0);
-
-	record_temp_fan(info, &ar, &(avalon->temp));
-	applog(LOG_INFO,
-	       "Avalon: Fan1: %d/m, Fan2: %d/m, Fan3: %d/m\t"
-	       "Temp1: %dC, Temp2: %dC, Temp3: %dC, TempMAX: %dC",
-	       info->fan0, info->fan1, info->fan2,
-	       info->temp0, info->temp1, info->temp2, info->temp_max);
-	info->temp_history_index++;
-	info->temp_sum += info->temp2;
-	applog(LOG_DEBUG, "Avalon: temp_index: %d, temp_count: %d, temp_old: %d",
-	       info->temp_history_index, info->temp_history_count, info->temp_old);
-	if (info->temp_history_index == info->temp_history_count) {
-		adjust_fan(info);
-		info->temp_history_index = 0;
-		info->temp_sum = 0;
-	}
-
-	/*
-	 * FIXME: Each work split to 10 pieces, each piece send to a
-	 * asic(256MHs). one work can be mulit-nonce back. it is not
-	 * easy calculate correct hash on such situation. so I simplely
-	 * add each nonce to hash_count. base on Utility/m hash_count*2
-	 * give a very good result.
-	 *
-	 * Any patch will be great.
-	 */
-	return (hash_count * 2);
-}
-
-// minerloop_scanhash hacked to handle Avalon's many processors
-static
-void minerloop_avalon(struct thr_info *mythr)
-{
-	const int thr_id = mythr->id;
-	struct cgpu_info *cgpu = mythr->cgpu;
-	struct device_drv *api = cgpu->drv;
-	struct timeval tv_start, tv_end;
-	struct timeval tv_hashes;
-	uint32_t max_nonce = api->can_limit_work ? api->can_limit_work(mythr) : 0xffffffff;
-	int64_t hashes;
-	struct avalon_info *info = avalon_info[cgpu->device_id];
-	int i;
-	int avalon_get_work_count = info->miner_count;
-	struct work **work = calloc(1,
-				    avalon_get_work_count * sizeof(struct work *));
-	if (!work)
-		quit(1, "Faile on Avalon calloc");
-	const bool primary = (!mythr->device_thread) || mythr->primary_thread;
-	
-	while (1) {
-		mythr->work_restart = false;
-		for (i = 0; i < avalon_get_work_count; i++)
-			request_work(mythr);
-		for (i = 0; i < avalon_get_work_count; i++) {
-			work[i] = get_work(mythr);
-			work[i]->blk.nonce = 0;
+	avalon_rotate_array(avalon);
+
+	if (hash_count) {
+		record_temp_fan(info, &ar, &(avalon->temp));
+		applog(LOG_INFO,
+		       "Avalon: Fan1: %d/m, Fan2: %d/m, Fan3: %d/m\t"
+		       "Temp1: %dC, Temp2: %dC, Temp3: %dC, TempMAX: %dC",
+		       info->fan0, info->fan1, info->fan2,
+		       info->temp0, info->temp1, info->temp2, info->temp_max);
+		info->temp_history_index++;
+		info->temp_sum += avalon->temp;
+		applog(LOG_DEBUG, "Avalon: temp_index: %d, temp_count: %d, temp_old: %d",
+		       info->temp_history_index, info->temp_history_count, info->temp_old);
+		if (info->temp_history_index == info->temp_history_count) {
+			adjust_fan(info);
+			info->temp_history_index = 0;
+			info->temp_sum = 0;
 		}
-		for (i = 0; i < avalon_get_work_count; i++) {
-			if (api->prepare_work && !api->prepare_work(mythr, work[i])) {
-				applog(LOG_ERR, "work prepare failed, exiting "
-					"mining thread %d", thr_id);
-				break;
-			}
-			gettimeofday(&(work[i]->tv_work_start), NULL);
-		}
-		
-		do {
-			thread_reportin(mythr);
-			gettimeofday(&tv_start, NULL);
-			hashes = api->scanhash_queue(mythr, work, max_nonce);
-			gettimeofday(&tv_end, NULL);
-			thread_reportin(mythr);
-			
-			timersub(&tv_end, &tv_start, &tv_hashes);
-			if (!hashes_done(mythr, hashes, &tv_hashes, api->can_limit_work ? &max_nonce : NULL))
-				goto disabled;
-			
-			if (unlikely(mythr->work_restart)) {
-				/* Apart from device_thread 0, we stagger the
-				 * starting of every next thread to try and get
-				 * all devices busy before worrying about
-				 * getting work for their extra threads */
-				if (!primary) {
-					struct timespec rgtp;
-
-					rgtp.tv_sec = 0;
-					rgtp.tv_nsec = 250 * mythr->device_thread * 1000000;
-					nanosleep(&rgtp, NULL);
-				}
-				break;
-			}
-			
-			if (unlikely(mythr->pause || cgpu->deven != DEV_ENABLED))
-disabled:
-				mt_disable(mythr);
-		} while (false);
 	}
-	free(work);
+
+	/* This hashmeter is just a utility counter based on returned shares */
+	return hash_count;
 }
 
-static struct api_data *avalon_drv_stats(struct cgpu_info *cgpu)
+static struct api_data *avalon_api_stats(struct cgpu_info *cgpu)
 {
 	struct api_data *root = NULL;
-	struct avalon_info *info = avalon_info[cgpu->device_id];
+	struct avalon_info *info = avalon_infos[cgpu->device_id];
+	int i;
 
 	root = api_add_int(root, "baud", &(info->baud), false);
 	root = api_add_int(root, "miner_count", &(info->miner_count),false);
 	root = api_add_int(root, "asic_count", &(info->asic_count), false);
-	root = api_add_int(root, "read_count", &(info->read_count), false);
 	root = api_add_int(root, "timeout", &(info->timeout), false);
 	root = api_add_int(root, "frequency", &(info->frequency), false);
 
@@ -1057,30 +989,12 @@ static struct api_data *avalon_drv_stats(struct cgpu_info *cgpu)
 	root = api_add_int(root, "temp_max", &(info->temp_max), false);
 
 	root = api_add_int(root, "no_matching_work", &(info->no_matching_work), false);
-	root = api_add_int(root, "matching_work_count1", &(info->matching_work[0]), false);
-	root = api_add_int(root, "matching_work_count2", &(info->matching_work[1]), false);
-	root = api_add_int(root, "matching_work_count3", &(info->matching_work[2]), false);
-	root = api_add_int(root, "matching_work_count4", &(info->matching_work[3]), false);
-	root = api_add_int(root, "matching_work_count5", &(info->matching_work[4]), false);
-	root = api_add_int(root, "matching_work_count6", &(info->matching_work[5]), false);
-	root = api_add_int(root, "matching_work_count7", &(info->matching_work[6]), false);
-	root = api_add_int(root, "matching_work_count8", &(info->matching_work[7]), false);
-	root = api_add_int(root, "matching_work_count9", &(info->matching_work[8]), false);
-	root = api_add_int(root, "matching_work_count10", &(info->matching_work[9]), false);
-	root = api_add_int(root, "matching_work_count11", &(info->matching_work[10]), false);
-	root = api_add_int(root, "matching_work_count12", &(info->matching_work[11]), false);
-	root = api_add_int(root, "matching_work_count13", &(info->matching_work[12]), false);
-	root = api_add_int(root, "matching_work_count14", &(info->matching_work[13]), false);
-	root = api_add_int(root, "matching_work_count15", &(info->matching_work[14]), false);
-	root = api_add_int(root, "matching_work_count16", &(info->matching_work[15]), false);
-	root = api_add_int(root, "matching_work_count17", &(info->matching_work[16]), false);
-	root = api_add_int(root, "matching_work_count18", &(info->matching_work[17]), false);
-	root = api_add_int(root, "matching_work_count19", &(info->matching_work[18]), false);
-	root = api_add_int(root, "matching_work_count20", &(info->matching_work[19]), false);
-	root = api_add_int(root, "matching_work_count21", &(info->matching_work[20]), false);
-	root = api_add_int(root, "matching_work_count22", &(info->matching_work[21]), false);
-	root = api_add_int(root, "matching_work_count23", &(info->matching_work[22]), false);
-	root = api_add_int(root, "matching_work_count24", &(info->matching_work[23]), false);
+	for (i = 0; i < info->miner_count; i++) {
+		char mcw[24];
+
+		sprintf(mcw, "match_work_count%d", i + 1);
+		root = api_add_int(root, mcw, &(info->matching_work[i]), false);
+	}
 
 	return root;
 }
@@ -1095,9 +1009,10 @@ struct device_drv avalon_drv = {
 	.name = "AVA",
 	.drv_detect = avalon_detect,
 	.thread_prepare = avalon_prepare,
-	.minerloop = minerloop_avalon,
-	.scanhash_queue = avalon_scanhash,
-	.get_api_stats = avalon_drv_stats,
+	.minerloop = hash_queued_work,
+	.queue_full = avalon_fill,
+	.scanwork = avalon_scanhash,
+	.get_api_stats = avalon_api_stats,
 	.reinit_device = avalon_init,
 	.thread_shutdown = avalon_shutdown,
 };

driver-avalon.h

@@ -10,7 +10,8 @@
 #ifndef AVALON_H
 #define AVALON_H
 
-#define AVALON_TIME_FACTOR 10
+#ifdef USE_AVALON
+
 #define AVALON_RESET_FAULT_DECISECONDS 1
 #define AVALON_MINER_THREADS 1
 
@@ -71,8 +72,6 @@ struct avalon_result {
 } __attribute__((packed, aligned(4)));
 
 struct avalon_info {
-	int read_count;
-
 	int baud;
 	int miner_count;
 	int asic_count;
@@ -94,16 +93,13 @@ struct avalon_info {
 
 	int no_matching_work;
 	int matching_work[AVALON_DEFAULT_MINER_NUM];
-	struct work *bulk0[AVALON_DEFAULT_MINER_NUM];
-	struct work *bulk1[AVALON_DEFAULT_MINER_NUM];
-	struct work *bulk2[AVALON_DEFAULT_MINER_NUM];
-	struct work *bulk3[AVALON_DEFAULT_MINER_NUM];
 
 	int frequency;
 };
 
 #define AVALON_WRITE_SIZE (sizeof(struct avalon_task))
 #define AVALON_READ_SIZE (sizeof(struct avalon_result))
+#define AVALON_ARRAY_SIZE 4
 
 #define AVA_GETS_ERROR -1
 #define AVA_GETS_OK 0
@@ -130,4 +126,5 @@ ASSERT1(sizeof(uint32_t) == 4);
 
 extern struct avalon_info **avalon_info;
 
+#endif /* USE_AVALON */
 #endif	/* AVALON_H */

driver-bitforce.c

@@ -22,6 +22,7 @@
 #include "deviceapi.h"
 #include "miner.h"
 #include "fpgautils.h"
+#include "util.h"
 
 #define BITFORCE_SLEEP_MS 500
 #define BITFORCE_TIMEOUT_S 7
@@ -286,7 +287,7 @@ static bool bitforce_thread_prepare(struct thr_info *thr)
 	bitforce->device_fd = fdDev;
 
 	applog(LOG_INFO, "%s: Opened %s", bitforce->dev_repr, bitforce->device_path);
-	gettimeofday(&now, NULL);
+	cgtime(&now);
 	get_datestamp(bitforce->init, &now);
 
 	return true;
@@ -699,7 +700,7 @@ re_send:
 	bitforce_cmd2(fdDev, data->xlink_id, pdevbuf, sizeof(pdevbuf), data->next_work_cmd, ob, data->next_work_obsz);
 	if (!pdevbuf[0] || !strncasecmp(pdevbuf, "B", 1)) {
 		mutex_unlock(mutexp);
-		gettimeofday(&tv_now, NULL);
+		cgtime(&tv_now);
 		timer_set_delay(&thr->tv_poll, &tv_now, WORK_CHECK_INTERVAL_MS * 1000);
 		data->poll_func = 1;
 		return;
@@ -727,7 +728,7 @@ re_send:
 
 	dbg_block_data(bitforce);
 
-	gettimeofday(&tv_now, NULL);
+	cgtime(&tv_now);
 	bitforce->work_start_tv = tv_now;
 	
 	timer_set_delay(&thr->tv_morework, &tv_now, bitforce->sleep_ms * 1000);
@@ -764,7 +765,7 @@ int bitforce_zox(struct thr_info *thr, const char *cmd)
 		size_t szleft = sizeof(data->noncebuf) - cls, sz;
 		
 		if (count && data->queued)
-			gettimeofday(&bitforce->work_start_tv, NULL);
+			cgtime(&bitforce->work_start_tv);
 		
 		while (true)
 		{
@@ -813,7 +814,7 @@ void bitforce_job_get_results(struct thr_info *thr, struct work *work)
 	bool stale;
 	int count;
 
-	gettimeofday(&now, NULL);
+	cgtime(&now);
 	timersub(&now, &bitforce->work_start_tv, &elapsed);
 	bitforce->wait_ms = tv_to_ms(elapsed);
 	bitforce->polling = true;
@@ -849,7 +850,7 @@ void bitforce_job_get_results(struct thr_info *thr, struct work *work)
 		const char *cmd = (data->proto == BFP_QUEUE) ? "ZOX" : "ZFX";
 		count = bitforce_zox(thr, cmd);
 
-		gettimeofday(&now, NULL);
+		cgtime(&now);
 		timersub(&now, &bitforce->work_start_tv, &elapsed);
 
 		if (elapsed.tv_sec >= BITFORCE_LONG_TIMEOUT_S) {
@@ -1381,7 +1382,7 @@ bool bitforce_send_queue(struct thr_info *thr)
 	}
 	
 	if (!data->queued)
-		gettimeofday(&data->tv_hashmeter_start, NULL);
+		cgtime(&data->tv_hashmeter_start);
 	
 	queued_ok = atoi(&buf[9]);
 	data->queued += queued_ok;
@@ -1511,7 +1512,7 @@ next_qline: (void)0;
 		applog(LOG_DEBUG, "%"PRIpreprv": Received %d queue results after %ums; Wait time unchanged (queued<=%d)",
 		       bitforce->proc_repr, count, bitforce->sleep_ms, data->queued);
 	
-	gettimeofday(&tv_now, NULL);
+	cgtime(&tv_now);
 	timersub(&tv_now, &data->tv_hashmeter_start, &tv_elapsed);
 	hashes_done(thr, (uint64_t)bitforce->nonces * count, &tv_elapsed, NULL);
 	data->tv_hashmeter_start = tv_now;

driver-cpu.c

@@ -75,7 +75,6 @@ static inline void affine_to_cpu(int __maybe_unused id, int __maybe_unused cpu)
 
 
 /* TODO: resolve externals */
-extern void submit_work_async(const struct work *work_in, struct timeval *tv);
 extern char *set_int_range(const char *arg, int *i, int min, int max);
 extern int dev_from_id(int thr_id);
 
@@ -840,7 +839,7 @@ CPUSearch:
 	/* if nonce found, submit work */
 	if (unlikely(rc)) {
 		applog(LOG_DEBUG, "%"PRIpreprv" found something?", thr->cgpu->proc_repr);
-		submit_work_async(work, NULL);
+		submit_nonce(thr, work, last_nonce);
 		work->blk.nonce = last_nonce + 1;
 		goto CPUSearch;
 	}

driver-icarus.c

@@ -242,7 +242,7 @@ int icarus_gets(unsigned char *buf, int fd, struct timeval *tv_finish, struct th
 			return ICA_GETS_ERROR;
 
 		if (first)
-			gettimeofday(tv_finish, NULL);
+			cgtime(tv_finish);
 
 		if (ret >= read_amount)
 		{
@@ -576,7 +576,7 @@ bool icarus_detect_custom(const char *devpath, struct device_drv *api, struct IC
 
 	hex2bin(ob_bin, golden_ob, sizeof(ob_bin));
 	icarus_write(fd, ob_bin, sizeof(ob_bin));
-	gettimeofday(&tv_start, NULL);
+	cgtime(&tv_start);
 
 	memset(nonce_bin, 0, sizeof(nonce_bin));
 	icarus_gets(nonce_bin, fd, &tv_finish, NULL, 1);
@@ -672,7 +672,7 @@ static bool icarus_prepare(struct thr_info *thr)
 	icarus->device_fd = fd;
 
 	applog(LOG_INFO, "Opened Icarus on %s", icarus->device_path);
-	gettimeofday(&now, NULL);
+	cgtime(&now);
 	get_datestamp(icarus->init, &now);
 
 	struct icarus_state *state;
@@ -766,7 +766,7 @@ static bool icarus_start_work(struct thr_info *thr, const unsigned char *ob_bin)
 	int ret;
 	char *ob_hex;
 
-	gettimeofday(&state->tv_workstart, NULL);
+	cgtime(&state->tv_workstart);
 
 	ret = icarus_write(fd, ob_bin, 64);
 	if (ret) {
@@ -996,7 +996,7 @@ static int64_t icarus_scanhash(struct thr_info *thr, struct work *work,
 	&&  !was_hw_error
 	&&  ((nonce & info->nonce_mask) > END_CONDITION)
 	&&  ((nonce & info->nonce_mask) < (info->nonce_mask & ~END_CONDITION))) {
-		gettimeofday(&tv_history_start, NULL);
+		cgtime(&tv_history_start);
 
 		history0 = &(info->history[0]);
 
@@ -1085,7 +1085,7 @@ static int64_t icarus_scanhash(struct thr_info *thr, struct work *work,
 					Hs, W, read_count, fullnonce);
 		}
 		info->history_count++;
-		gettimeofday(&tv_history_finish, NULL);
+		cgtime(&tv_history_finish);
 
 		timersub(&tv_history_finish, &tv_history_start, &tv_history_finish);
 		timeradd(&tv_history_finish, &(info->history_time), &(info->history_time));

driver-opencl.c

@@ -44,6 +44,7 @@
 #include "findnonce.h"
 #include "ocl.h"
 #include "adl.h"
+#include "util.h"
 
 /* TODO: cleanup externals ********************/
 
@@ -1355,7 +1356,7 @@ select_cgpu:
 
 		thr->rolling = thr->cgpu->rolling = 0;
 		/* Reports the last time we tried to revive a sick GPU */
-		gettimeofday(&thr->sick, NULL);
+		cgtime(&thr->sick);
 		if (!pthread_cancel(thr->pth)) {
 			applog(LOG_WARNING, "Thread %d still exists, killing it off", thr_id);
 		} else
@@ -1395,7 +1396,7 @@ select_cgpu:
 		applog(LOG_WARNING, "Thread %d restarted", thr_id);
 	}
 
-	gettimeofday(&now, NULL);
+	cgtime(&now);
 	get_datestamp(sel_cgpu->init, &now);
 
 	proc_enable(cgpu);
@@ -1587,9 +1588,10 @@ static bool opencl_thread_prepare(struct thr_info *thr)
 	int virtual_gpu = cgpu->virtual_gpu;
 	int i = thr->id;
 	static bool failmessage = false;
+	int buffersize = opt_scrypt ? SCRYPT_BUFFERSIZE : BUFFERSIZE;
 
 	if (!blank_res)
-		blank_res = calloc(BUFFERSIZE, 1);
+		blank_res = calloc(buffersize, 1);
 	if (!blank_res) {
 		applog(LOG_ERR, "Failed to calloc in opencl_thread_init");
 		return false;
@@ -1651,7 +1653,7 @@ static bool opencl_thread_prepare(struct thr_info *thr)
 		}
 	}
 	applog(LOG_INFO, "initCl() finished. Found %s", name);
-	gettimeofday(&now, NULL);
+	cgtime(&now);
 	get_datestamp(cgpu->init, &now);
 
 	have_opencl = true;
@@ -1668,6 +1670,7 @@ static bool opencl_thread_init(struct thr_info *thr)
 	cl_int status = 0;
 	thrdata = calloc(1, sizeof(*thrdata));
 	thr->cgpu_data = thrdata;
+	int buffersize = opt_scrypt ? SCRYPT_BUFFERSIZE : BUFFERSIZE;
 
 	if (!thrdata) {
 		applog(LOG_ERR, "Failed to calloc in opencl_thread_init");
@@ -1695,7 +1698,7 @@ static bool opencl_thread_init(struct thr_info *thr)
 			break;
 	}
 
-	thrdata->res = calloc(BUFFERSIZE, 1);
+	thrdata->res = calloc(buffersize, 1);
 
 	if (!thrdata->res) {
 		free(thrdata);
@@ -1704,7 +1707,7 @@ static bool opencl_thread_init(struct thr_info *thr)
 	}
 
 	status |= clEnqueueWriteBuffer(clState->commandQueue, clState->outputBuffer, CL_TRUE, 0,
-			BUFFERSIZE, blank_res, 0, NULL, NULL);
+				       buffersize, blank_res, 0, NULL, NULL);
 	if (unlikely(status != CL_SUCCESS)) {
 		applog(LOG_ERR, "Error: clEnqueueWriteBuffer failed.");
 		return false;
@@ -1745,13 +1748,15 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
 	size_t globalThreads[1];
 	size_t localThreads[1] = { clState->wsize };
 	int64_t hashes;
+	int found = opt_scrypt ? SCRYPT_FOUND : FOUND;
+	int buffersize = opt_scrypt ? SCRYPT_BUFFERSIZE : BUFFERSIZE;
 
 	/* Windows' timer resolution is only 15ms so oversample 5x */
 	if (gpu->dynamic && (++gpu->intervals * dynamic_us) > 70000) {
 		struct timeval tv_gpuend;
 		double gpu_us;
 
-		gettimeofday(&tv_gpuend, NULL);
+		cgtime(&tv_gpuend);
 		gpu_us = us_tdiff(&tv_gpuend, &gpu->tv_gpustart) / gpu->intervals;
 		if (gpu_us > dynamic_us) {
 			if (gpu->intensity > MIN_INTENSITY)
@@ -1789,7 +1794,7 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
 	}
 
 	status = clEnqueueReadBuffer(clState->commandQueue, clState->outputBuffer, CL_FALSE, 0,
-			BUFFERSIZE, thrdata->res, 0, NULL, NULL);
+				     buffersize, thrdata->res, 0, NULL, NULL);
 	if (unlikely(status != CL_SUCCESS)) {
 		applog(LOG_ERR, "Error: clEnqueueReadBuffer failed error %d. (clEnqueueReadBuffer)", status);
 		return -1;
@@ -1804,17 +1809,17 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
 	clFinish(clState->commandQueue);
 
 	/* FOUND entry is used as a counter to say how many nonces exist */
-	if (thrdata->res[FOUND]) {
+	if (thrdata->res[found]) {
 		/* Clear the buffer again */
 		status = clEnqueueWriteBuffer(clState->commandQueue, clState->outputBuffer, CL_FALSE, 0,
-				BUFFERSIZE, blank_res, 0, NULL, NULL);
+					      buffersize, blank_res, 0, NULL, NULL);
 		if (unlikely(status != CL_SUCCESS)) {
 			applog(LOG_ERR, "Error: clEnqueueWriteBuffer failed.");
 			return -1;
 		}
 		applog(LOG_DEBUG, "GPU %d found something?", gpu->device_id);
 		postcalc_hash_async(thr, work, thrdata->res);
-		memset(thrdata->res, 0, BUFFERSIZE);
+		memset(thrdata->res, 0, buffersize);
 		/* This finish flushes the writebuffer set with CL_FALSE in clEnqueueWriteBuffer */
 		clFinish(clState->commandQueue);
 	}

driver-ztex.c

@@ -33,6 +33,7 @@
 #include "dynclock.h"
 #include "fpgautils.h"
 #include "libztex.h"
+#include "util.h"
 
 #define GOLDEN_BACKLOG 5
 
@@ -357,7 +358,7 @@ static bool ztex_prepare(struct thr_info *thr)
 	struct cgpu_info *cgpu = thr->cgpu;
 	struct libztex_device *ztex = cgpu->device_ztex;
 
-	gettimeofday(&now, NULL);
+	cgtime(&now);
 	get_datestamp(cgpu->init, &now);
 
 	if (cgpu->proc_id)

findnonce.c

@@ -136,7 +136,7 @@ void precalc_hash(dev_blk_ctx *blk, uint32_t *state, uint32_t *data)
 struct pc_data {
 	struct thr_info *thr;
 	struct work work;
-	uint32_t res[MAXBUFFERS];
+	uint32_t res[SCRYPT_MAXBUFFERS];
 	pthread_t pth;
 	int found;
 };
@@ -146,21 +146,22 @@ static void *postcalc_hash(void *userdata)
 	struct pc_data *pcd = (struct pc_data *)userdata;
 	struct thr_info *thr = pcd->thr;
 	unsigned int entry = 0;
+	int found = opt_scrypt ? SCRYPT_FOUND : FOUND;
 
 	pthread_detach(pthread_self());
 	RenameThread("postcalchsh");
 
 	/* To prevent corrupt values in FOUND from trying to read beyond the
 	 * end of the res[] array */
-	if (unlikely(pcd->res[FOUND] & ~FOUND)) {
+	if (unlikely(pcd->res[found] & ~found)) {
 		applog(LOG_WARNING, "%"PRIpreprv": invalid nonce count - HW error",
 				thr->cgpu->proc_repr);
 		hw_errors++;
 		thr->cgpu->hw_errors++;
-		pcd->res[FOUND] &= FOUND;
+		pcd->res[found] &= found;
 	}
 
-	for (entry = 0; entry < pcd->res[FOUND]; entry++) {
+	for (entry = 0; entry < pcd->res[found]; entry++) {
 		uint32_t nonce = pcd->res[entry];
 
 		applog(LOG_DEBUG, "OCL NONCE %u found in slot %d", nonce, entry);
@@ -176,6 +177,8 @@ static void *postcalc_hash(void *userdata)
 void postcalc_hash_async(struct thr_info *thr, struct work *work, uint32_t *res)
 {
 	struct pc_data *pcd = malloc(sizeof(struct pc_data));
+	int buffersize;
+
 	if (unlikely(!pcd)) {
 		applog(LOG_ERR, "Failed to malloc pc_data in postcalc_hash_async");
 		return;
@@ -185,7 +188,8 @@ void postcalc_hash_async(struct thr_info *thr, struct work *work, uint32_t *res)
 		.thr = thr,
 	};
 	__copy_work(&pcd->work, work);
-	memcpy(&pcd->res, res, BUFFERSIZE);
+	buffersize = opt_scrypt ? SCRYPT_BUFFERSIZE : BUFFERSIZE;
+	memcpy(&pcd->res, res, buffersize);
 
 	if (pthread_create(&pcd->pth, NULL, postcalc_hash, (void *)pcd)) {
 		applog(LOG_ERR, "Failed to create postcalc_hash thread");

findnonce.h

@@ -8,6 +8,10 @@
 #define BUFFERSIZE (sizeof(uint32_t) * MAXBUFFERS)
 #define FOUND (0x0F)
 
+#define SCRYPT_MAXBUFFERS (0x100)
+#define SCRYPT_BUFFERSIZE (sizeof(uint32_t) * SCRYPT_MAXBUFFERS)
+#define SCRYPT_FOUND (0xFF)
+
 #ifdef HAVE_OPENCL
 extern void precalc_hash(dev_blk_ctx *blk, uint32_t *state, uint32_t *data);
 extern void postcalc_hash_async(struct thr_info *thr, struct work *work, uint32_t *res);

fpgautils.c

@@ -1,5 +1,6 @@
 /*
  * Copyright 2012-2013 Luke Dashjr
+ * Copyright 2013 Con Kolivas
  * Copyright 2012 Andrew Smith
  * Copyright 2013 Xiangfu
  *
@@ -27,8 +28,8 @@
 #ifndef WIN32
 #include <errno.h>
 #include <termios.h>
-#include <sys/stat.h>
 #include <sys/ioctl.h>
+#include <sys/stat.h>
 #include <unistd.h>
 #include <fcntl.h>
 #ifndef O_CLOEXEC
@@ -57,10 +58,12 @@ enum {
 
 #ifdef HAVE_LIBUDEV
 #include <libudev.h>
+#include <sys/ioctl.h>
 #endif
 
 #include "elist.h"
 #include "logging.h"
+#include "miner.h"
 #include "fpgautils.h"
 
 #define SEARCH_NEEDLES_BEGIN()  {  \
@@ -928,5 +931,19 @@ int set_serial_rts(int fd, int rts)
 	ioctl(fd, TIOCMSET, &flags);
 	return flags & TIOCM_CTS;
 }
+#else
+int get_serial_cts(const int fd)
+{
+	if (!fd)
+		return -1;
+	const HANDLE fh = (HANDLE)_get_osfhandle(fd);
+	if (!fh)
+		return -1;
 
+	DWORD flags;
+	if (!GetCommModemStatus(fh, &flags))
+		return -1;
+
+	return (flags & MS_CTS_ON) ? 1 : 0;
+}
 #endif // ! WIN32

logging.c

@@ -15,6 +15,7 @@
 #include "compat.h"
 #include "logging.h"
 #include "miner.h"
+// #include "util.h"
 
 bool opt_debug = false;
 bool opt_debug_console = false;  // Only used if opt_debug is also enabled
@@ -99,7 +100,7 @@ static void log_generic(int prio, const char *fmt, va_list ap)
 		struct tm _tm;
 		struct tm *tm = &_tm;
 
-		gettimeofday(&tv, NULL);
+		cgtime(&tv);
 
 		localtime_r(&tv.tv_sec, tm);
 
@@ -119,6 +120,7 @@ static void log_generic(int prio, const char *fmt, va_list ap)
 
 			va_copy(apc, ap);
 			vfprintf(stderr, f, apc);	/* atomic write to stderr */
+			va_end(apc);
 			fflush(stderr);
 		}
 

miner.h

@@ -289,9 +289,6 @@ struct device_drv {
 	uint64_t (*can_limit_work)(struct thr_info *);
 	bool (*thread_init)(struct thr_info *);
 	bool (*prepare_work)(struct thr_info *, struct work *);
-#ifdef USE_AVALON
-	int64_t (*scanhash_queue)(struct thr_info *, struct work **, int64_t);
-#endif
 	int64_t (*scanhash)(struct thr_info *, struct work *, int64_t);
 	int64_t (*scanwork)(struct thr_info *);
 
@@ -436,6 +433,12 @@ struct cgpu_info {
 		struct ft232r_device_handle *device_ft232r;
 #endif
 	};
+#ifdef USE_AVALON
+	struct work **works;
+	int work_array;
+	int queued;
+	int results;
+#endif
 #ifdef USE_BITFORCE
 	struct timeval work_start_tv;
 	unsigned int wait_ms;
@@ -535,6 +538,7 @@ struct cgpu_info {
 
 	pthread_rwlock_t qlock;
 	struct work *queued_work;
+	unsigned int queued_count;
 };
 
 extern void renumber_cgpu(struct cgpu_info *);
@@ -603,13 +607,6 @@ struct thr_info {
 	notifier_t work_restart_notifier;
 };
 
-extern int thr_info_create(struct thr_info *thr, pthread_attr_t *attr, void *(*start) (void *), void *arg);
-extern void thr_info_cancel(struct thr_info *thr);
-extern void thr_info_freeze(struct thr_info *thr);
-extern void nmsleep(unsigned int msecs);
-extern double us_tdiff(struct timeval *end, struct timeval *start);
-extern double tdiff(struct timeval *end, struct timeval *start);
-
 struct string_elist {
 	char *string;
 	bool free_me;
@@ -839,7 +836,9 @@ extern bool opt_restart;
 extern char *opt_icarus_options;
 extern char *opt_icarus_timing;
 extern bool opt_worktime;
+#ifdef USE_AVALON
 extern char *opt_avalon_options;
+#endif
 #ifdef USE_BITFORCE
 extern bool opt_bfl_noncerange;
 #endif
@@ -870,6 +869,7 @@ extern int opt_scantime;
 extern int opt_expiry;
 
 extern cglock_t control_lock;
+extern pthread_mutex_t stats_lock;
 extern pthread_mutex_t hash_lock;
 extern pthread_mutex_t console_lock;
 extern cglock_t ch_lock;
@@ -882,6 +882,7 @@ extern void clear_stratum_shares(struct pool *pool);
 extern void hashmeter2(struct thr_info *);
 extern bool stale_work(struct work *, bool share);
 extern bool stale_work_future(struct work *, bool share, unsigned long ustime);
+extern void set_target(unsigned char *dest_target, double diff);
 
 extern void kill_work(void);
 extern void app_restart(void);
@@ -1159,6 +1160,8 @@ struct work {
 	unsigned char	target[32];
 	unsigned char	hash[32];
 
+	uint64_t	share_diff;
+
 	int		rolls;
 
 	dev_blk_ctx	blk;
@@ -1191,6 +1194,10 @@ struct work {
 	
 	double		work_difficulty;
 
+	// Allow devices to identify work if multiple sub-devices
+	// DEPRECATED: New code should be using multiple processors instead
+	char		subid;
+
 	blktemplate_t	*tmpl;
 	int		*tmpl_refcount;
 	unsigned int	dataid;
@@ -1208,10 +1215,11 @@ struct work {
 };
 
 extern void get_datestamp(char *, struct timeval *);
+extern void inc_hw_errors(struct thr_info *thr);
 enum test_nonce2_result {
-	TNR_GOOD,
-	TNR_HIGH,
-	TNR_BAD,
+	TNR_GOOD = 1,
+	TNR_HIGH = 0,
+	TNR_BAD = -1,
 };
 extern enum test_nonce2_result _test_nonce2(struct work *, uint32_t nonce, bool checktarget);
 #define test_nonce(work, nonce, checktarget)  (_test_nonce2(work, nonce, checktarget) == TNR_GOOD)

ocl.c

@@ -212,7 +212,12 @@ CL_API_ENTRY cl_int CL_API_CALL
                        cl_event *       /* event */) CL_API_SUFFIX__VERSION_1_0;
 
 int opt_platform_id = -1;
+#ifdef __APPLE__
+// Apple OpenCL doesn't like using binaries this way
+bool opt_opencl_binaries;
+#else
 bool opt_opencl_binaries = true;
+#endif
 
 char *file_contents(const char *filename, int *length)
 {
@@ -668,6 +673,8 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
 
 	if (cgpu->work_size && cgpu->work_size <= clState->max_work_size)
 		clState->wsize = cgpu->work_size;
+	else if (opt_scrypt)
+		clState->wsize = 256;
 	else if (strstr(name, "Tahiti"))
 		clState->wsize = 64;
 	else
@@ -1027,7 +1034,8 @@ built:
 			applog(LOG_ERR, "Error %d: clCreateBuffer (CLbuffer0)", status);
 			return NULL;
 		}
-	}
+		clState->outputBuffer = clCreateBuffer(clState->context, CL_MEM_WRITE_ONLY, SCRYPT_BUFFERSIZE, NULL, &status);
+	} else
 #endif
 	clState->outputBuffer = clCreateBuffer(clState->context, CL_MEM_WRITE_ONLY, BUFFERSIZE, NULL, &status);
 	if (status != CL_SUCCESS) {

scrypt.c

@@ -405,7 +405,7 @@ static void scrypt_1024_1_1_256_sp(const uint32_t* input, char* scratchpad, uint
 /* 131583 rounded up to 4 byte alignment */
 #define SCRATCHBUF_SIZE	(131584)
 
-void scrypt_outputhash(struct work *work)
+void scrypt_regenhash(struct work *work)
 {
 	uint32_t data[20];
 	char *scratchbuf;
@@ -419,10 +419,12 @@ void scrypt_outputhash(struct work *work)
 	flip32(ohash, ohash);
 }
 
+static const uint32_t diff1targ = 0x0000ffff;
+
 /* Used externally as confirmation of correct OCL code */
-bool scrypt_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t nonce)
+int scrypt_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t nonce)
 {
-	uint32_t tmp_hash7, Htarg = ((const uint32_t *)ptarget)[7];
+	uint32_t tmp_hash7, Htarg = le32toh(((const uint32_t *)ptarget)[7]);
 	uint32_t data[20], ohash[8];
 	char *scratchbuf;
 
@@ -432,7 +434,15 @@ bool scrypt_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t no
 	scrypt_1024_1_1_256_sp(data, scratchbuf, ohash);
 	tmp_hash7 = be32toh(ohash[7]);
 
-	return (tmp_hash7 <= Htarg);
+	applog(LOG_DEBUG, "htarget %08lx diff1 %08lx hash %08lx",
+				(long unsigned int)Htarg,
+				(long unsigned int)diff1targ,
+				(long unsigned int)tmp_hash7);
+	if (tmp_hash7 > diff1targ)
+		return -1;
+	if (tmp_hash7 > Htarg)
+		return 0;
+	return 1;
 }
 
 bool scanhash_scrypt(struct thr_info *thr, const unsigned char __maybe_unused *pmidstate,
@@ -444,7 +454,7 @@ bool scanhash_scrypt(struct thr_info *thr, const unsigned char __maybe_unused *p
 	char *scratchbuf;
 	uint32_t data[20];
 	uint32_t tmp_hash7;
-	uint32_t Htarg = ((const uint32_t *)ptarget)[7];
+	uint32_t Htarg = le32toh(((const uint32_t *)ptarget)[7]);
 	bool ret = false;
 
 	be32enc_vect(data, (const uint32_t *)pdata, 19);
@@ -459,7 +469,7 @@ bool scanhash_scrypt(struct thr_info *thr, const unsigned char __maybe_unused *p
 		uint32_t ostate[8];
 
 		*nonce = ++n;
-		data[19] = n;
+		data[19] = htobe32(n);
 		scrypt_1024_1_1_256_sp(data, scratchbuf, ostate);
 		tmp_hash7 = be32toh(ostate[7]);
 

scrypt.h

@@ -4,19 +4,19 @@
 #include "miner.h"
 
 #ifdef USE_SCRYPT
-extern bool scrypt_test(unsigned char *pdata, const unsigned char *ptarget,
+extern int scrypt_test(unsigned char *pdata, const unsigned char *ptarget,
 			uint32_t nonce);
-extern void scrypt_outputhash(struct work *work);
+extern void scrypt_regenhash(struct work *work);
 
 #else /* USE_SCRYPT */
-static inline bool scrypt_test(__maybe_unused unsigned char *pdata,
+static inline int scrypt_test(__maybe_unused unsigned char *pdata,
 			       __maybe_unused const unsigned char *ptarget,
 			       __maybe_unused uint32_t nonce)
 {
-	return false;
+	return 0;
 }
 
-static inline void scrypt_outputhash(__maybe_unused struct work *work)
+static inline void scrypt_regenhash(__maybe_unused struct work *work)
 {
 }
 #endif /* USE_SCRYPT */

scrypt130302.cl → scrypt130511.cl

@@ -808,8 +808,8 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 	unshittify(X);
 }
 
-#define FOUND (0x0F)
-#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
+#define SCRYPT_FOUND (0xFF)
+#define SETFOUND(Xnonce) output[output[SCRYPT_FOUND]++] = Xnonce
 
 __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
 __kernel void search(__global const uint4 * restrict input,

util.c

@@ -32,6 +32,7 @@
 # include <pthread_np.h>
 #endif
 #ifndef WIN32
+#include <fcntl.h>
 # ifdef __linux
 #  include <sys/prctl.h>
 # endif
@@ -40,9 +41,11 @@
 # include <netinet/tcp.h>
 # include <netdb.h>
 #else
+# include <windows.h>
 # include <winsock2.h>
 # include <mstcpip.h>
 # include <ws2tcpip.h>
+# include <mmsystem.h>
 #endif
 
 #include "miner.h"
@@ -234,21 +237,35 @@ out:
 
 static int keep_sockalive(SOCKETTYPE fd)
 {
+	const int tcp_one = 1;
 	const int tcp_keepidle = 45;
 	const int tcp_keepintvl = 30;
-	const int keepalive = 1;
 	int ret = 0;
 
+	if (unlikely(setsockopt(fd, SOL_SOCKET, SO_KEEPALIVE, (const char *)&tcp_one, sizeof(tcp_one))))
+		ret = 1;
 
 #ifndef WIN32
-	const int tcp_keepcnt = 1;
+	int flags = fcntl(fd, F_GETFL, 0);
 
-	if (unlikely(setsockopt(fd, SOL_SOCKET, SO_KEEPALIVE, &keepalive, sizeof(keepalive))))
-		ret = 1;
+	fcntl(fd, F_SETFL, O_NONBLOCK | flags);
+#else
+	u_long flags = 1;
 
-# ifdef __linux
+	ioctlsocket(fd, FIONBIO, &flags);
+#endif
+
+	if (!opt_delaynet)
+#ifndef __linux
+		if (unlikely(setsockopt(fd, IPPROTO_TCP, TCP_NODELAY, (const void *)&tcp_one, sizeof(tcp_one))))
+#else /* __linux */
+		if (unlikely(setsockopt(fd, SOL_TCP, TCP_NODELAY, (const void *)&tcp_one, sizeof(tcp_one))))
+#endif /* __linux */
+			ret = 1;
+
+#ifdef __linux
 
-	if (unlikely(setsockopt(fd, SOL_TCP, TCP_KEEPCNT, &tcp_keepcnt, sizeof(tcp_keepcnt))))
+	if (unlikely(setsockopt(fd, SOL_TCP, TCP_KEEPCNT, &tcp_one, sizeof(tcp_one))))
 		ret = 1;
 
 	if (unlikely(setsockopt(fd, SOL_TCP, TCP_KEEPIDLE, &tcp_keepidle, sizeof(tcp_keepidle))))
@@ -256,15 +273,16 @@ static int keep_sockalive(SOCKETTYPE fd)
 
 	if (unlikely(setsockopt(fd, SOL_TCP, TCP_KEEPINTVL, &tcp_keepintvl, sizeof(tcp_keepintvl))))
 		ret = 1;
-# endif /* __linux */
-# ifdef __APPLE_CC__
+#endif /* __linux */
+
+#ifdef __APPLE_CC__
 
 	if (unlikely(setsockopt(fd, IPPROTO_TCP, TCP_KEEPALIVE, &tcp_keepintvl, sizeof(tcp_keepintvl))))
 		ret = 1;
 
-# endif /* __APPLE_CC__ */
+#endif /* __APPLE_CC__ */
 
-#else /* WIN32 */
+#ifdef WIN32
 
 	const int zero = 0;
 	struct tcp_keepalive vals;
@@ -274,9 +292,6 @@ static int keep_sockalive(SOCKETTYPE fd)
 
 	DWORD outputBytes;
 
-	if (unlikely(setsockopt(fd, SOL_SOCKET, SO_KEEPALIVE, (const char *)&keepalive, sizeof(keepalive))))
-		ret = 1;
-
 	if (unlikely(WSAIoctl(fd, SIO_KEEPALIVE_VALS, &vals, sizeof(vals), NULL, 0, &outputBytes, NULL, NULL)))
 		ret = 1;
 
@@ -307,7 +322,7 @@ static void last_nettime(struct timeval *last)
 static void set_nettime(void)
 {
 	wr_lock(&netacc_lock);
-	gettimeofday(&nettime, NULL);
+	cgtime(&nettime);
 	wr_unlock(&netacc_lock);
 }
 
@@ -463,7 +478,7 @@ void json_rpc_call_async(CURL *curl, const char *url,
 			long long now_msecs, last_msecs;
 			struct timeval now, last;
 
-			gettimeofday(&now, NULL);
+			cgtime(&now);
 			last_nettime(&last);
 			now_msecs = (long long)now.tv_sec * 1000;
 			now_msecs += now.tv_usec / 1000;
@@ -775,7 +790,7 @@ bool hash_target_check_v(const unsigned char *hash, const unsigned char *target)
 		applog(LOG_DEBUG, " Proof: %s\nTarget: %s\nTrgVal? %s",
 			hash_str,
 			target_str,
-			rc ? "YES (hash < target)" :
+			rc ? "YES (hash <= target)" :
 			     "no (false positive; hash > target)");
 
 		free(hash_str);
@@ -831,9 +846,7 @@ void tq_free(struct thread_q *tq)
 static void tq_freezethaw(struct thread_q *tq, bool frozen)
 {
 	mutex_lock(&tq->mutex);
-
 	tq->frozen = frozen;
-
 	pthread_cond_signal(&tq->cond);
 	mutex_unlock(&tq->mutex);
 }
@@ -861,14 +874,12 @@ bool tq_push(struct thread_q *tq, void *data)
 	INIT_LIST_HEAD(&ent->q_node);
 
 	mutex_lock(&tq->mutex);
-
 	if (!tq->frozen) {
 		list_add_tail(&ent->q_node, &tq->q);
 	} else {
 		free(ent);
 		rc = false;
 	}
-
 	pthread_cond_signal(&tq->cond);
 	mutex_unlock(&tq->mutex);
 
@@ -882,7 +893,6 @@ void *tq_pop(struct thread_q *tq, const struct timespec *abstime)
 	int rc;
 
 	mutex_lock(&tq->mutex);
-
 	if (!list_empty(&tq->q))
 		goto pop;
 
@@ -894,16 +904,15 @@ void *tq_pop(struct thread_q *tq, const struct timespec *abstime)
 		goto out;
 	if (list_empty(&tq->q))
 		goto out;
-
 pop:
 	ent = list_entry(tq->q.next, struct tq_ent, q_node);
 	rval = ent->data;
 
 	list_del(&ent->q_node);
 	free(ent);
-
 out:
 	mutex_unlock(&tq->mutex);
+
 	return rval;
 }
 
@@ -1050,6 +1059,9 @@ void nmsleep(unsigned int msecs)
 	int ret;
 	ldiv_t d;
 
+#ifdef WIN32
+	timeBeginPeriod(1);
+#endif
 	d = ldiv(msecs, 1000);
 	tleft.tv_sec = d.quot;
 	tleft.tv_nsec = d.rem * 1000000;
@@ -1058,6 +1070,48 @@ void nmsleep(unsigned int msecs)
 		twait.tv_nsec = tleft.tv_nsec;
 		ret = nanosleep(&twait, &tleft);
 	} while (ret == -1 && errno == EINTR);
+#ifdef WIN32
+	timeEndPeriod(1);
+#endif
+}
+
+/* This is a cgminer gettimeofday wrapper. Since we always call gettimeofday
+ * with tz set to NULL, and windows' default resolution is only 15ms, this
+ * gives us higher resolution times on windows. */
+void cgtime(struct timeval *tv)
+{
+#ifdef WIN32
+	timeBeginPeriod(1);
+#endif
+	gettimeofday(tv, NULL);
+#ifdef WIN32
+	timeEndPeriod(1);
+#endif
+}
+
+void subtime(struct timeval *a, struct timeval *b)
+{
+	timersub(a, b, b);
+}
+
+void addtime(struct timeval *a, struct timeval *b)
+{
+	timeradd(a, b, b);
+}
+
+bool time_more(struct timeval *a, struct timeval *b)
+{
+	return timercmp(a, b, >);
+}
+
+bool time_less(struct timeval *a, struct timeval *b)
+{
+	return timercmp(a, b, <);
+}
+
+void copy_time(struct timeval *dest, const struct timeval *src)
+{
+	memcpy(dest, src, sizeof(struct timeval));
 }
 
 /* Returns the microseconds difference between end and start times as a double */
@@ -1136,7 +1190,7 @@ static enum send_ret __stratum_send(struct pool *pool, char *s, ssize_t len)
 	len++;
 
 	while (len > 0 ) {
-		struct timeval timeout = {0, 0};
+		struct timeval timeout = {1, 0};
 		ssize_t sent;
 		fd_set wd;
 
@@ -1144,7 +1198,13 @@ static enum send_ret __stratum_send(struct pool *pool, char *s, ssize_t len)
 		FD_SET(sock, &wd);
 		if (select(sock + 1, NULL, &wd, NULL, &timeout) < 1)
 			return SEND_SELECTFAIL;
+#ifdef __APPLE__
+		sent = send(pool->sock, s + ssent, len, SO_NOSIGPIPE);
+#elif WIN32
 		sent = send(pool->sock, s + ssent, len, 0);
+#else
+		sent = send(pool->sock, s + ssent, len, MSG_NOSIGNAL);
+#endif
 		if (sent < 0) {
 			if (!sock_blocks())
 				return SEND_SENDFAIL;
@@ -1203,7 +1263,7 @@ static bool socket_full(struct pool *pool, bool wait)
 	if (wait)
 		timeout.tv_sec = 60;
 	else
-		timeout.tv_sec = 0;
+		timeout.tv_sec = 1;
 	if (select(sock + 1, &rd, NULL, NULL, &timeout) > 0)
 		return true;
 	return false;
@@ -1248,7 +1308,8 @@ static void recalloc_sock(struct pool *pool, size_t len)
 	if (new < pool->sockbuf_size)
 		return;
 	new = new + (RBUFSIZE - (new % RBUFSIZE));
-	applog(LOG_DEBUG, "Recallocing pool sockbuf to %lu", (unsigned long)new);
+	// Avoid potentially recursive locking
+	// applog(LOG_DEBUG, "Recallocing pool sockbuf to %lu", (unsigned long)new);
 	pool->sockbuf = realloc(pool->sockbuf, new);
 	if (!pool->sockbuf)
 		quit(1, "Failed to realloc pool sockbuf in recalloc_sock");
@@ -1273,7 +1334,7 @@ char *recv_line(struct pool *pool)
 		enum recv_ret ret = RECV_OK;
 		struct timeval rstart, now;
 
-		gettimeofday(&rstart, NULL);
+		cgtime(&rstart);
 		if (!socket_full(pool, true)) {
 			applog(LOG_DEBUG, "Timed out waiting for data on socket_full");
 			goto out;
@@ -1292,7 +1353,7 @@ char *recv_line(struct pool *pool)
 				break;
 			}
 			if (n < 0) {
-				if (!sock_blocks()) {
+				if (!sock_blocks() || !socket_full(pool, false)) {
 					ret = RECV_RECVFAIL;
 					break;
 				}
@@ -1301,7 +1362,7 @@ char *recv_line(struct pool *pool)
 				recalloc_sock(pool, slen);
 				strcat(pool->sockbuf, s);
 			}
-			gettimeofday(&now, NULL);
+			cgtime(&now);
 		} while (tdiff(&now, &rstart) < 60 && !strstr(pool->sockbuf, "\n"));
 		mutex_unlock(&pool->stratum_lock);
 
@@ -1798,6 +1859,7 @@ static bool setup_stratum_curl(struct pool *pool)
 	if (pool->sockbuf)
 		pool->sockbuf[0] = '\0';
 	mutex_unlock(&pool->stratum_lock);
+
 	curl = pool->stratum_curl;
 
 	if (!pool->sockbuf) {
@@ -1815,7 +1877,8 @@ static bool setup_stratum_curl(struct pool *pool)
 	curl_easy_setopt(curl, CURLOPT_ERRORBUFFER, curl_err_str);
 	curl_easy_setopt(curl, CURLOPT_NOSIGNAL, 1);
 	curl_easy_setopt(curl, CURLOPT_URL, s);
-	curl_easy_setopt(curl, CURLOPT_TCP_NODELAY, 1);
+	if (!opt_delaynet)
+		curl_easy_setopt(curl, CURLOPT_TCP_NODELAY, 1);
 
 	/* We use DEBUGFUNCTION to count bytes sent/received, and verbose is needed
 	 * to enable it */
@@ -1889,6 +1952,7 @@ void suspend_stratum(struct pool *pool)
 {
 	clear_sockbuf(pool);
 	applog(LOG_INFO, "Closing socket for stratum pool %d", pool->pool_no);
+
 	mutex_lock(&pool->stratum_lock);
 	pool->stratum_active = pool->stratum_notify = false;
 	if (pool->stratum_curl) {

util.h

@@ -56,6 +56,7 @@
 #endif
 extern char *json_dumps_ANY(json_t *, size_t flags);
 
+struct thr_info;
 struct pool;
 enum dev_reason;
 struct cgpu_info;
@@ -71,6 +72,18 @@ extern void real_block_target(unsigned char *target, const unsigned char *data);
 extern bool hash_target_check(const unsigned char *hash, const unsigned char *target);
 extern bool hash_target_check_v(const unsigned char *hash, const unsigned char *target);
 
+int thr_info_create(struct thr_info *thr, pthread_attr_t *attr, void *(*start) (void *), void *arg);
+void thr_info_freeze(struct thr_info *thr);
+void thr_info_cancel(struct thr_info *thr);
+void nmsleep(unsigned int msecs);
+void cgtime(struct timeval *tv);
+void subtime(struct timeval *a, struct timeval *b);
+void addtime(struct timeval *a, struct timeval *b);
+bool time_more(struct timeval *a, struct timeval *b);
+bool time_less(struct timeval *a, struct timeval *b);
+void copy_time(struct timeval *dest, const struct timeval *src);
+double us_tdiff(struct timeval *end, struct timeval *start);
+double tdiff(struct timeval *end, struct timeval *start);
 bool stratum_send(struct pool *pool, char *s, ssize_t len);
 bool sock_full(struct pool *pool);
 char *recv_line(struct pool *pool);