Merge branch 'master' into workaround_icarus_uart_issue

Conflicts:
	driver-icarus.c

Makefile.am

@@ -9,11 +9,12 @@ endif
 
 EXTRA_DIST	= example.conf m4/gnulib-cache.m4 linux-usb-cgminer \
 		  ADL_SDK/readme.txt api-example.php miner.php	\
-		  API.class API.java api-example.c windows-build.txt
+		  API.class API.java api-example.c windows-build.txt \
+		  bitstreams/*
 
 SUBDIRS		= lib compat ccan
 
-INCLUDES	= $(PTHREAD_FLAGS) -fno-strict-aliasing $(JANSSON_INCLUDES)
+INCLUDES	= $(PTHREAD_FLAGS) -fno-strict-aliasing $(JANSSON_INCLUDES) $(USB_FLAGS)
 
 bin_PROGRAMS	= cgminer
 
@@ -22,7 +23,7 @@ bin_SCRIPTS	= *.cl
 cgminer_LDFLAGS	= $(PTHREAD_FLAGS)
 cgminer_LDADD	= $(DLOPEN_FLAGS) @LIBCURL_LIBS@ @JANSSON_LIBS@ @PTHREAD_LIBS@ \
 		  @OPENCL_LIBS@ @NCURSES_LIBS@ @PDCURSES_LIBS@ @WS2_LIBS@ \
-		  @UDEV_LIBS@ \
+		  @UDEV_LIBS@ @USB_LIBS@ \
 		  @MATH_LIBS@ lib/libgnu.a ccan/libccan.a
 cgminer_CPPFLAGS = -I$(top_builddir)/lib -I$(top_srcdir)/lib @OPENCL_FLAGS@
 
@@ -37,7 +38,7 @@ cgminer_SOURCES	+= logging.c
 
 # GPU sources, TODO: make them selectable
 # the GPU portion extracted from original main.c
-cgminer_SOURCES += device-gpu.h device-gpu.c
+cgminer_SOURCES += driver-opencl.h driver-opencl.c
 
 # the original GPU related sources, unchanged
 cgminer_SOURCES += ocl.c ocl.h findnonce.c findnonce.h
@@ -53,7 +54,7 @@ cgminer_SOURCES	+= \
 		  sha256_altivec_4way.c
 
 # the CPU portion extracted from original main.c
-cgminer_SOURCES += device-cpu.h device-cpu.c
+cgminer_SOURCES += driver-cpu.h driver-cpu.c
 
 if HAS_YASM
 AM_CFLAGS	= -DHAS_YASM
@@ -68,9 +69,15 @@ endif # HAS_YASM
 endif # HAS_CPUMINE
 
 if HAS_BITFORCE
-cgminer_SOURCES += bitforce.c
+cgminer_SOURCES += driver-bitforce.c
 endif
 
 if HAS_ICARUS
-cgminer_SOURCES += icarus.c
+cgminer_SOURCES += driver-icarus.c
+endif
+
+if HAS_ZTEX
+cgminer_SOURCES += driver-ztex.c libztex.c
+bitstreamsdir = $(bindir)/bitstreams
+dist_bitstreams_DATA = bitstreams/*
 endif

NEWS

@@ -1,3 +1,27 @@
+Version 2.3.3 - April 15, 2012
+
+- Don't even display that cpumining is disabled on ./configure to discourage
+people from enabling it.
+- Do a complete cgminer restart if the ATI Display Library fails, as it does on
+windows after running for some time, when fanspeed reporting fails.
+- Cache the initial arguments passed to cgminer and implement an attempted
+restart option from the settings menu.
+- Disable per-device status lines when there are more than 8 devices since
+screen output will be corrupted, enumerating them to the log output instead at
+startup.
+- Reuse Vals[] array more than W[] till they're re-initialised on the second
+sha256 cycle in poclbm kernel.
+- Minor variable alignment in poclbm kernel.
+- Make sure to disable devices with any status not being DEV_ENABLED to ensure
+that thermal cutoff code works as it was setting the status to DEV_RECOVER.
+- Re-initialising ADL simply made the driver fail since it is corruption over
+time within the windows driver that's responsible. Revert "Attempt to
+re-initialise ADL should a device that previously reported fanspeed stops
+reporting it."
+- Microoptimise poclbm kernel by ordering Val variables according to usage
+frequency.
+
+
 Version 2.3.2 - March 31, 2012
 
 - Damping small changes in hashrate so dramatically has the tendency to always

README

@@ -596,6 +596,15 @@ An example request in both formats to set GPU 0 fan to 80%:
 The format of each reply (unless stated otherwise) is a STATUS section
 followed by an optional detail section
 
+From API verion 1.7 onwards, reply strings in JSON and Text have the
+necessary escaping as required to avoid ambiguity - they didn't before 1.7
+For JSON the 2 characters '"' and '\' are escaped with a '\' before them
+For Text the 4 characters '|' ',' '=' and '\' are escaped the same way
+
+Only user entered information will contain characters that require being
+escaped, such as Pool URL, User and Password or the Config save filename,
+when they are returned in messages or as their values by the API
+
 For API version 1.4 and later:
 
 The STATUS section is:
@@ -622,7 +631,7 @@ The STATUS section is:
    This defaults to the cgminer version but is the value of --api-description
    if it was specified at runtime.
 
-For API version 1.6:
+For API version 1.7:
 
 The list of requests - a (*) means it requires privileged access - and replies are:
 
@@ -700,6 +709,12 @@ The list of requests - a (*) means it requires privileged access - and replies a
                               stating the results of disabling pool N
                               The Msg includes the pool URL
 
+ removepool|N (*)
+               none           There is no reply section just the STATUS section
+                              stating the results of removing pool N
+                              The Msg includes the pool URL
+                              N.B. all details for the pool will be lost
+
  gpuenable|N (*)
                none           There is no reply section just the STATUS section
                               stating the results of the enable request
@@ -794,9 +809,8 @@ api-example.c - a 'C' program to access the API (with source code)
 
 miner.php - an example web page to access the API
  This includes buttons and inputs to attempt access to the privileged commands
- You must modify the 2 lines near the top to change where it looks for cgminer
-  $miner = '127.0.0.1'; # hostname or IP address
-  $port = 4028;
+ Read the top of the file (miner.php) for details of how to tune the display
+ and also to use the option to display a multi-rig summary
 
 ---
 
@@ -873,21 +887,14 @@ any further.
 
 Q: Can you change the autofan/autogpu to change speeds in a different manner?
 A: The defaults are sane and safe. I'm not interested in changing them
-further. The starting fan speed is set to 85% in auto-fan mode as a safety
-precaution, but if a specific fan speed has been set, it will use that first
-before adjusting automatically.
-
-Q: The fanspeed starts at 85% with --auto-fan. Can I set it lower?
-A: The initial fanspeed will always start at 85% unless you choose your own
-value with --gpu-fan. In this case it will use the value you give it with
---gpu-fan as the first fanspeed, but it will also use this as the maximum fan
-speed unless overheat is detected.
+further. The starting fan speed is set to 50% in auto-fan mode as a safety
+precaution.
 
 Q: Why is my efficiency above/below 100%?
 A: Efficiency simply means how many shares you return for the amount of work
 you request. It does not correlate with efficient use of your hardware, and is
 a measure of a combination of hardware speed, block luck, pool design and other
-factors.
+factors
 
 Q: What are the best parameters to pass for X pool/hardware/device.
 A: Virtually always, the DEFAULT parameters give the best results. Most user
@@ -912,7 +919,7 @@ this time.
 
 Q: Which ATI SDK is the best for cgminer?
 A: At the moment, versions 2.4 and 2.5 work the best. If you are forced to use
-the 2.6 SDK, -v 1 might help, along with not decreasing your memory clock speed.
+the 2.6 SDK.
 
 Q: I have multiple SDKs installed, can I choose which one it uses?
 A: Run cgminer with the -n option and it will list all the platforms currently

adl.c

@@ -352,6 +352,9 @@ void init_adl(int nDevs)
 		}
 
 		applog(LOG_INFO, "GPU %d %s hardware monitoring enabled", gpu, lpInfo[i].strAdapterName);
+		if (gpus[gpu].name)
+			free(gpus[gpu].name);
+		gpus[gpu].name = lpInfo[i].strAdapterName;
 		gpus[gpu].has_adl = true;
 		/* Flag adl as active if any card is successfully activated */
 		adl_active = true;
@@ -639,8 +642,6 @@ static inline int __gpu_fanspeed(struct gpu_adl *ga)
 	return ga->lpFanSpeedValue.iFanSpeed;
 }
 
-static void reinit_adl(void);
-
 int gpu_fanspeed(int gpu)
 {
 	struct gpu_adl *ga;
@@ -669,8 +670,6 @@ static int __gpu_fanpercent(struct gpu_adl *ga)
 	return ga->lpFanSpeedValue.iFanSpeed;
 }
 
-
-
 int gpu_fanpercent(int gpu)
 {
 	struct gpu_adl *ga;
@@ -684,9 +683,14 @@ int gpu_fanpercent(int gpu)
 	ret = __gpu_fanpercent(ga);
 	unlock_adl();
 	if (unlikely(ga->has_fanspeed && ret == -1)) {
-		applog(LOG_WARNING, "GPU %d stopped reporting fanspeed", gpu);
-		applog(LOG_WARNING, "Will attempt to re-initialise ADL");
-		reinit_adl();
+		applog(LOG_WARNING, "GPU %d stopped reporting fanspeed due to driver corruption", gpu);
+		if (opt_restart) {
+			applog(LOG_WARNING, "Restart enabled, will restart cgminer");
+			applog(LOG_WARNING, "You can disable this with the --no-restart option");
+			app_restart();
+		}
+		applog(LOG_WARNING, "Disabling fanspeed monitoring on this device");
+		ga->has_fanspeed = false;
 	}
 	return ret;
 }
@@ -1368,7 +1372,7 @@ void clear_adl(int nDevs)
 	free_adl();
 }
 
-static void reinit_adl(void)
+void reinit_adl(void)
 {
 	bool ret;
 	lock_adl();
@@ -1377,8 +1381,7 @@ static void reinit_adl(void)
 	if (!ret) {
 		adl_active = false;
 		applog(LOG_WARNING, "Attempt to re-initialise ADL has failed, disabling");
-	} else
-		applog(LOG_WARNING, "ADL re-initialisation complete");
+	}
 	unlock_adl();
 }
 #endif /* HAVE_ADL */

adl.h

@@ -19,10 +19,12 @@ bool gpu_stats(int gpu, float *temp, int *engineclock, int *memclock, float *vdd
 void change_gpusettings(int gpu);
 void gpu_autotune(int gpu, enum dev_enable *denable);
 void clear_adl(int nDevs);
+void reinit_adl(void);
 #else /* HAVE_ADL */
 #define adl_active (0)
 static inline void init_adl(int nDevs) {}
 static inline void change_gpusettings(int gpu) { }
 static inline void clear_adl(int nDevs) {}
+static inline void reinit_adl(void) {}
 #endif
 #endif

api.c

@@ -25,7 +25,7 @@
 
 #include "compat.h"
 #include "miner.h"
-#include "device-cpu.h" /* for algo_names[], TODO: re-factor dependency */
+#include "driver-cpu.h" /* for algo_names[], TODO: re-factor dependency */
 
 #if defined(unix) || defined(__APPLE__)
 	#include <errno.h>
@@ -157,7 +157,7 @@ static const char *COMMA = ",";
 static const char SEPARATOR = '|';
 static const char GPUSEP = ',';
 
-static const char *APIVERSION = "1.6";
+static const char *APIVERSION = "1.7";
 static const char *DEAD = "Dead";
 static const char *SICK = "Sick";
 static const char *NOSTART = "NoStart";
@@ -337,6 +337,10 @@ static const char *JSON_PARAMETER = "parameter";
 #define MSG_PGAUNW 65
 #endif
 
+#define MSG_REMLASTP 66
+#define MSG_ACTPOOL 67
+#define MSG_REMPOOL 68
+
 enum code_severity {
 	SEVERITY_ERR,
 	SEVERITY_WARN,
@@ -456,6 +460,9 @@ struct CODES {
  { SEVERITY_ERR,   MSG_INVPDP,	PARAM_STR,	"Invalid addpool details '%s'" },
  { SEVERITY_ERR,   MSG_TOOMANYP,PARAM_NONE,	"Reached maximum number of pools (%d)" },
  { SEVERITY_SUCC,  MSG_ADDPOOL,	PARAM_STR,	"Added pool '%s'" },
+ { SEVERITY_ERR,   MSG_REMLASTP,PARAM_POOL,	"Cannot remove last pool %d:'%s'" },
+ { SEVERITY_ERR,   MSG_ACTPOOL, PARAM_POOL,	"Cannot remove active pool %d:'%s'" },
+ { SEVERITY_SUCC,  MSG_REMPOOL, PARAM_BOTH,	"Removed pool %d:'%s'" },
  { SEVERITY_SUCC,  MSG_NOTIFY,	PARAM_NONE,	"Notify" },
  { SEVERITY_FAIL, 0, 0, NULL }
 };
@@ -483,6 +490,68 @@ extern struct device_api bitforce_api;
 extern struct device_api icarus_api;
 #endif
 
+// This is only called when expected to be needed (rarely)
+// i.e. strings outside of the codes control (input from the user)
+static char *escape_string(char *str, bool isjson)
+{
+	char *buf, *ptr;
+	int count;
+
+	count = 0;
+	for (ptr = str; *ptr; ptr++) {
+		switch (*ptr) {
+		case ',':
+		case '|':
+		case '=':
+			if (!isjson)
+				count++;
+			break;
+		case '"':
+			if (isjson)
+				count++;
+			break;
+		case '\\':
+			count++;
+			break;
+		}
+	}
+
+	if (count == 0)
+		return str;
+
+	buf = malloc(strlen(str) + count + 1);
+	if (unlikely(!buf))
+		quit(1, "Failed to malloc escape buf");
+
+	ptr = buf;
+	while (*str)
+		switch (*str) {
+		case ',':
+		case '|':
+		case '=':
+			if (!isjson)
+				*(ptr++) = '\\';
+			*(ptr++) = *(str++);
+			break;
+		case '"':
+			if (isjson)
+				*(ptr++) = '\\';
+			*(ptr++) = *(str++);
+			break;
+		case '\\':
+			*(ptr++) = '\\';
+			*(ptr++) = *(str++);
+			break;
+		default:
+			*(ptr++) = *(str++);
+			break;
+		}
+
+	*ptr = '\0';
+
+	return buf;
+}
+
 #if defined(USE_BITFORCE) || defined(USE_ICARUS)
 static int numpgas()
 {
@@ -1102,6 +1171,8 @@ static void poolstatus(__maybe_unused SOCKETTYPE c, __maybe_unused char *param,
 {
 	char buf[BUFSIZ];
 	char *status, *lp;
+	char *rpc_url;
+	char *rpc_user;
 	int i;
 
 	if (total_pools == 0) {
@@ -1133,27 +1204,40 @@ static void poolstatus(__maybe_unused SOCKETTYPE c, __maybe_unused char *param,
 		else
 			lp = (char *)NO;
 
+		rpc_url = escape_string(pool->rpc_url, isjson);
+		rpc_user = escape_string(pool->rpc_user, isjson);
+
 		if (isjson)
-			sprintf(buf, "%s{\"POOL\":%d,\"URL\":\"%s\",\"Status\":\"%s\",\"Priority\":%d,\"Long Poll\":\"%s\",\"Getworks\":%d,\"Accepted\":%d,\"Rejected\":%d,\"Discarded\":%d,\"Stale\":%d,\"Get Failures\":%d,\"Remote Failures\":%d}",
+			sprintf(buf, "%s{\"POOL\":%d,\"URL\":\"%s\",\"Status\":\"%s\",\"Priority\":%d,\"Long Poll\":\"%s\",\"Getworks\":%d,\"Accepted\":%d,\"Rejected\":%d,\"Discarded\":%d,\"Stale\":%d,\"Get Failures\":%d,\"Remote Failures\":%d,\"User\":\"%s\"}",
 				(i > 0) ? COMMA : "",
-				i, pool->rpc_url, status, pool->prio, lp,
+				i, rpc_url, status, pool->prio, lp,
 				pool->getwork_requested,
 				pool->accepted, pool->rejected,
 				pool->discarded_work,
 				pool->stale_shares,
 				pool->getfail_occasions,
-				pool->remotefail_occasions);
+				pool->remotefail_occasions,
+				rpc_user);
 		else
-			sprintf(buf, "POOL=%d,URL=%s,Status=%s,Priority=%d,Long Poll=%s,Getworks=%d,Accepted=%d,Rejected=%d,Discarded=%d,Stale=%d,Get Failures=%d,Remote Failures=%d%c",
-				i, pool->rpc_url, status, pool->prio, lp,
+			sprintf(buf, "POOL=%d,URL=%s,Status=%s,Priority=%d,Long Poll=%s,Getworks=%d,Accepted=%d,Rejected=%d,Discarded=%d,Stale=%d,Get Failures=%d,Remote Failures=%d,User=%s%c",
+				i, rpc_url, status, pool->prio, lp,
 				pool->getwork_requested,
 				pool->accepted, pool->rejected,
 				pool->discarded_work,
 				pool->stale_shares,
 				pool->getfail_occasions,
-				pool->remotefail_occasions, SEPARATOR);
+				pool->remotefail_occasions,
+				rpc_user, SEPARATOR);
 
 		strcat(io_buffer, buf);
+
+		if (rpc_url != pool->rpc_url)
+			free(rpc_url);
+		rpc_url = NULL;
+
+		if (rpc_user != pool->rpc_user)
+			free(rpc_user);
+		rpc_user = NULL;
 	}
 
 	if (isjson)
@@ -1443,6 +1527,7 @@ exitsama:
 static void addpool(__maybe_unused SOCKETTYPE c, char *param, bool isjson)
 {
 	char *url, *user, *pass;
+	char *ptr;
 
 	if (param == NULL || *param == '\0') {
 		strcpy(io_buffer, message(MSG_MISPDP, 0, NULL, isjson));
@@ -1450,7 +1535,11 @@ static void addpool(__maybe_unused SOCKETTYPE c, char *param, bool isjson)
 	}
 
 	if (!pooldetails(param, &url, &user, &pass)) {
-		strcpy(io_buffer, message(MSG_INVPDP, 0, param, isjson));
+		ptr = escape_string(param, isjson);
+		strcpy(io_buffer, message(MSG_INVPDP, 0, ptr, isjson));
+		if (ptr != param)
+			free(ptr);
+		ptr = NULL;
 		return;
 	}
 
@@ -1459,7 +1548,11 @@ static void addpool(__maybe_unused SOCKETTYPE c, char *param, bool isjson)
 		return;
 	}
 
-	strcpy(io_buffer, message(MSG_ADDPOOL, 0, url, isjson));
+	ptr = escape_string(url, isjson);
+	strcpy(io_buffer, message(MSG_ADDPOOL, 0, ptr, isjson));
+	if (ptr != url)
+		free(ptr);
+	ptr = NULL;
 }
 
 static void enablepool(__maybe_unused SOCKETTYPE c, char *param, bool isjson)
@@ -1535,6 +1628,57 @@ static void disablepool(__maybe_unused SOCKETTYPE c, char *param, bool isjson)
 	strcpy(io_buffer, message(MSG_DISPOOL, id, NULL, isjson));
 }
 
+static void removepool(__maybe_unused SOCKETTYPE c, char *param, bool isjson)
+{
+	struct pool *pool;
+	char *rpc_url;
+	bool dofree = false;
+	int id;
+
+	if (total_pools == 0) {
+		strcpy(io_buffer, message(MSG_NOPOOL, 0, NULL, isjson));
+		return;
+	}
+
+	if (param == NULL || *param == '\0') {
+		strcpy(io_buffer, message(MSG_MISPID, 0, NULL, isjson));
+		return;
+	}
+
+	id = atoi(param);
+	if (id < 0 || id >= total_pools) {
+		strcpy(io_buffer, message(MSG_INVPID, id, NULL, isjson));
+		return;
+	}
+
+	if (total_pools <= 1) {
+		strcpy(io_buffer, message(MSG_REMLASTP, id, NULL, isjson));
+		return;
+	}
+
+	pool = pools[id];
+	if (pool == current_pool())
+		switch_pools(NULL);
+
+	if (pool == current_pool()) {
+		strcpy(io_buffer, message(MSG_ACTPOOL, id, NULL, isjson));
+		return;
+	}
+
+	pool->enabled = false;
+	rpc_url = escape_string(pool->rpc_url, isjson);
+	if (rpc_url != pool->rpc_url)
+		dofree = true;
+
+	remove_pool(pool);
+
+	strcpy(io_buffer, message(MSG_REMPOOL, id, rpc_url, isjson));
+
+	if (dofree)
+		free(rpc_url);
+	rpc_url = NULL;
+}
+
 static bool splitgpuvalue(char *param, int *gpu, char **value, bool isjson)
 {
 	int id;
@@ -1792,6 +1936,7 @@ static void notify(__maybe_unused SOCKETTYPE c, __maybe_unused char *param, bool
 void dosave(__maybe_unused SOCKETTYPE c, char *param, bool isjson)
 {
 	FILE *fcfg;
+	char *ptr;
 
 	if (param == NULL || *param == '\0') {
 		strcpy(io_buffer, message(MSG_MISFN, 0, NULL, isjson));
@@ -1800,14 +1945,22 @@ void dosave(__maybe_unused SOCKETTYPE c, char *param, bool isjson)
 
 	fcfg = fopen(param, "w");
 	if (!fcfg) {
-		strcpy(io_buffer, message(MSG_BADFN, 0, param, isjson));
+		ptr = escape_string(param, isjson);
+		strcpy(io_buffer, message(MSG_BADFN, 0, ptr, isjson));
+		if (ptr != param)
+			free(ptr);
+		ptr = NULL;
 		return;
 	}
 
 	write_config(fcfg);
 	fclose(fcfg);
 
-	strcpy(io_buffer, message(MSG_SAVED, 0, param, isjson));
+	ptr = escape_string(param, isjson);
+	strcpy(io_buffer, message(MSG_SAVED, 0, ptr, isjson));
+	if (ptr != param)
+		free(ptr);
+	ptr = NULL;
 }
 
 struct CMDS {
@@ -1839,6 +1992,7 @@ struct CMDS {
 	{ "addpool",		addpool,	true },
 	{ "enablepool",		enablepool,	true },
 	{ "disablepool",	disablepool,	true },
+	{ "removepool",		removepool,	true },
 	{ "gpuintensity",	gpuintensity,	true },
 	{ "gpumem",		gpumem,		true },
 	{ "gpuengine",		gpuengine,	true },
@@ -1875,7 +2029,7 @@ static void send_result(SOCKETTYPE c, bool isjson)
 
 }
 
-static void tidyup()
+static void tidyup(void *arg)
 {
 	bye = 1;
 
@@ -2028,6 +2182,7 @@ void api(int api_thr_id)
 	bool did;
 	int i;
 
+	pthread_cleanup_push(tidyup, NULL);
 	my_thr_id = api_thr_id;
 
 	/* This should be done first to ensure curl has already called WSAStartup() in windows */
@@ -2238,5 +2393,5 @@ void api(int api_thr_id)
 		CLOSESOCKET(c);
 	}
 die:
-	tidyup();
+	pthread_cleanup_pop(true);
 }

bitstreams/LICENSE.txt

@@ -0,0 +1,24 @@
+All the bitstream files included in this directory that follow the name pattern ztex_*.bit are:
+
+----
+
+Copyright (C) 2009-2011 ZTEX GmbH.
+http://www.ztex.de
+
+
+This program is free software; you can redistribute it and/or modify
+it under the terms of the GNU General Public License version 3 as
+published by the Free Software Foundation.
+
+This program is distributed in the hope that it will be useful, but
+WITHOUT ANY WARRANTY; without even the implied warranty of
+MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+General Public License for more details.
+
+You should have received a copy of the GNU General Public License
+along with this program; if not, see http://www.gnu.org/licenses/.
+
+----
+
+You can find the original sources at the BTCMiner project home page: http://www.ztex.de/btcminer/
+

cgminer.c

@@ -44,8 +44,8 @@
 #include "miner.h"
 #include "findnonce.h"
 #include "adl.h"
-#include "device-cpu.h"
-#include "device-gpu.h"
+#include "driver-cpu.h"
+#include "driver-opencl.h"
 #include "bench_block.h"
 
 #if defined(unix)
@@ -54,7 +54,6 @@
 	#include <sys/wait.h>
 #endif
 
-
 enum workio_commands {
 	WC_GET_WORK,
 	WC_SUBMIT_WORK,
@@ -1016,6 +1015,9 @@ static char *opt_verusage_and_exit(const char *extra)
 #endif
 #ifdef USE_ICARUS
 		"icarus "
+#endif
+#ifdef USE_ZTEX
+		"ztex "
 #endif
 		"mining support.\n"
 		, packagename);
@@ -1162,7 +1164,10 @@ WINDOW *mainwin, *statuswin, *logwin;
 #endif
 double total_secs = 0.1;
 static char statusline[256];
+/* logstart is where the log window should start */
 static int devcursor, logstart, logcursor;
+/* statusy is where the status window goes up to in cases where it won't fit at startup */
+static int statusy;
 struct cgpu_info gpus[MAX_GPUDEVICES]; /* Maximum number apparently possible */
 struct cgpu_info *cpus;
 
@@ -1203,6 +1208,8 @@ static void get_statline(char *buf, struct cgpu_info *cgpu)
 	sprintf(buf, "%s%d ", cgpu->api->name, cgpu->device_id);
 	if (cgpu->api->get_statline_before)
 		cgpu->api->get_statline_before(buf, cgpu);
+	else
+		tailsprintf(buf, "               | ");
 	tailsprintf(buf, "(%ds):%.1f (avg):%.1f Mh/s | A:%d R:%d HW:%d U:%.2f/m",
 		opt_log_interval,
 		cgpu->rolling,
@@ -1255,7 +1262,7 @@ static void curses_print_status(void)
 	wclrtoeol(statuswin);
 	mvwprintw(statuswin, 5, 0, " Block: %s...  Started: %s", current_hash, blocktime);
 	mvwhline(statuswin, 6, 0, '-', 80);
-	mvwhline(statuswin, logstart - 1, 0, '-', 80);
+	mvwhline(statuswin, statusy - 1, 0, '-', 80);
 	mvwprintw(statuswin, devcursor - 1, 1, "[P]ool management %s[S]ettings [D]isplay options [Q]uit",
 		have_opencl ? "[G]PU management " : "");
 }
@@ -1272,28 +1279,34 @@ static void curses_print_devstatus(int thr_id)
 	struct cgpu_info *cgpu = thr_info[thr_id].cgpu;
 	char logline[255];
 
-		cgpu->utility = cgpu->accepted / ( total_secs ? total_secs : 1 ) * 60;
+	cgpu->utility = cgpu->accepted / ( total_secs ? total_secs : 1 ) * 60;
 
-	mvwprintw(statuswin, devcursor + cgpu->cgminer_id, 0, " %s %d: ", cgpu->api->name, cgpu->device_id);
+	/* Check this isn't out of the window size */
+	if (wmove(statuswin,devcursor + cgpu->cgminer_id, 0) == ERR)
+		return;
+	wprintw(statuswin, " %s %d: ", cgpu->api->name, cgpu->device_id);
 	if (cgpu->api->get_statline_before) {
 		logline[0] = '\0';
 		cgpu->api->get_statline_before(logline, cgpu);
 		wprintw(statuswin, "%s", logline);
 	}
-		if (cgpu->status == LIFE_DEAD)
-			wprintw(statuswin, "DEAD ");
-		else if (cgpu->status == LIFE_SICK)
-			wprintw(statuswin, "SICK ");
+	else
+		wprintw(statuswin, "               | ");
+
+	if (cgpu->status == LIFE_DEAD)
+		wprintw(statuswin, "DEAD ");
+	else if (cgpu->status == LIFE_SICK)
+		wprintw(statuswin, "SICK ");
 	else if (cgpu->deven == DEV_DISABLED)
 		wprintw(statuswin, "OFF  ");
 	else if (cgpu->deven == DEV_RECOVER)
 		wprintw(statuswin, "REST  ");
 	else
 		wprintw(statuswin, "%5.1f", cgpu->rolling);
-		adj_width(cgpu->accepted, &awidth);
-		adj_width(cgpu->rejected, &rwidth);
-		adj_width(cgpu->hw_errors, &hwwidth);
-		adj_width(cgpu->utility, &uwidth);
+	adj_width(cgpu->accepted, &awidth);
+	adj_width(cgpu->rejected, &rwidth);
+	adj_width(cgpu->hw_errors, &hwwidth);
+	adj_width(cgpu->utility, &uwidth);
 	wprintw(statuswin, "/%5.1fMh/s | A:%*d R:%*d HW:%*d U:%*.2f/m",
 			cgpu->total_mhashes / total_secs,
 			awidth, cgpu->accepted,
@@ -1307,7 +1320,7 @@ static void curses_print_devstatus(int thr_id)
 		wprintw(statuswin, "%s", logline);
 	}
 
-		wclrtoeol(statuswin);
+	wclrtoeol(statuswin);
 }
 #endif
 
@@ -1322,16 +1335,31 @@ static void print_status(int thr_id)
 static inline bool change_logwinsize(void)
 {
 	int x, y, logx, logy;
+	bool ret = false;
 
 	getmaxyx(mainwin, y, x);
-	getmaxyx(logwin, logy, logx);
+	if (x < 80 || y < 25)
+		return ret;
+
+	if (y > statusy + 2 && statusy < logstart) {
+		if (y - 2 < logstart)
+			statusy = y - 2;
+		else
+			statusy = logstart;
+		logcursor = statusy + 1;
+		mvwin(logwin, logcursor, 0);
+		wresize(statuswin, statusy, x);
+		ret = true;
+	}
+
 	y -= logcursor;
+	getmaxyx(logwin, logy, logx);
 	/* Detect screen size change */
-	if ((x != logx || y != logy) && x >= 80 && y >= 25) {
+	if (x != logx || y != logy) {
 		wresize(logwin, y, x);
-		return true;
+		ret = true;
 	}
-	return false;
+	return ret;
 }
 
 static void check_winsizes(void)
@@ -1342,7 +1370,12 @@ static void check_winsizes(void)
 		int y, x;
 
 		x = getmaxx(statuswin);
-		wresize(statuswin, logstart, x);
+		if (logstart > LINES - 2)
+			statusy = LINES - 2;
+		else
+			statusy = logstart;
+		logcursor = statusy + 1;
+		wresize(statuswin, statusy, x);
 		getmaxyx(mainwin, y, x);
 		y -= logcursor;
 		wresize(logwin, y, x);
@@ -1523,6 +1556,8 @@ static bool submit_upstream_work(const struct work *work)
 	res = json_object_get(val, "result");
 
 	if (!QUIET) {
+#ifndef MIPSEB
+// This one segfaults on my router for some reason
 		isblock = regeneratehash(work);
 		if (isblock)
 			found_blocks++;
@@ -1530,6 +1565,7 @@ static bool submit_upstream_work(const struct work *work)
 		sprintf(hashshow, "%08lx.%08lx.%08lx%s",
 			(unsigned long)(hash32[7]), (unsigned long)(hash32[6]), (unsigned long)(hash32[5]),
 			isblock ? " BLOCK!" : "");
+#endif
 	}
 
 	/* Theoretically threads could race when modifying accepted and
@@ -1773,8 +1809,7 @@ static void disable_curses(void)
 
 static void print_summary(void);
 
-/* This should be the common exit path */
-void kill_work(void)
+static void __kill_work(void)
 {
 	struct thr_info *thr;
 	int i;
@@ -1821,11 +1856,37 @@ void kill_work(void)
 	applog(LOG_DEBUG, "Killing off API thread");
 	thr = &thr_info[api_thr_id];
 	thr_info_cancel(thr);
+}
+
+/* This should be the common exit path */
+void kill_work(void)
+{
+	__kill_work();
 
 	quit(0, "Shutdown signal received.");
 }
 
-void quit(int status, const char *format, ...);
+static char **initial_args;
+
+static void clean_up(void);
+
+void app_restart(void)
+{
+	applog(LOG_WARNING, "Attempting to restart %s", packagename);
+
+	__kill_work();
+	clean_up();
+
+#if defined(unix)
+	if (forkpid > 0) {
+		kill(forkpid, SIGTERM);
+		forkpid = 0;
+	}
+#endif
+
+	execv(initial_args[0], initial_args);
+	applog(LOG_WARNING, "Failed to restart application");
+}
 
 static void sighandler(int __maybe_unused sig)
 {
@@ -2355,10 +2416,11 @@ static void display_pool_summary(struct pool *pool)
 		unlock_curses();
 	}
 }
+#endif
 
 /* We can't remove the memory used for this struct pool because there may
  * still be work referencing it. We just remove it from the pools list */
-static void remove_pool(struct pool *pool)
+void remove_pool(struct pool *pool)
 {
 	int i, last_pool = total_pools - 1;
 	struct pool *other;
@@ -2379,7 +2441,6 @@ static void remove_pool(struct pool *pool)
 	pool->pool_no = total_pools;
 	total_pools--;
 }
-#endif
 
 void write_config(FILE *fcfg)
 {
@@ -2751,7 +2812,8 @@ static void set_options(void)
 	clear_logwin();
 retry:
 	wlogprint("\n[L]ongpoll: %s\n", want_longpoll ? "On" : "Off");
-	wlogprint("[Q]ueue: %d\n[S]cantime: %d\n[E]xpiry: %d\n[R]etries: %d\n[P]ause: %d\n[W]rite config file\n",
+	wlogprint("[Q]ueue: %d\n[S]cantime: %d\n[E]xpiry: %d\n[R]etries: %d\n"
+		  "[P]ause: %d\n[W]rite config file\n[C]gminer restart\n",
 		opt_queue, opt_scantime, opt_expiry, opt_retries, opt_fail_pause);
 	wlogprint("Select an option or any other key to return\n");
 	input = getch();
@@ -2844,6 +2906,13 @@ retry:
 		fclose(fcfg);
 		goto retry;
 
+	} else if (!strncasecmp(&input, "c", 1)) {
+		wlogprint("Are you sure?\n");
+		input = getch();
+		if (!strncasecmp(&input, "y", 1))
+			app_restart();
+		else
+			clear_logwin();
 	} else
 		clear_logwin();
 
@@ -3600,7 +3669,7 @@ void *miner_thread(void *userdata)
 				tv_lastupdate = tv_end;
 			}
 
-			if (unlikely(mythr->pause || cgpu->deven == DEV_DISABLED)) {
+			if (unlikely(mythr->pause || cgpu->deven != DEV_ENABLED)) {
 				applog(LOG_WARNING, "Thread %d being disabled", thr_id);
 				mythr->rolling = mythr->cgpu->rolling = 0;
 				applog(LOG_DEBUG, "Popping wakeup ping in miner thread");
@@ -4007,8 +4076,8 @@ static void log_print_status(struct cgpu_info *cgpu)
 {
 	char logline[255];
 
-		get_statline(logline, cgpu);
-		applog(LOG_WARNING, "%s", logline);
+	get_statline(logline, cgpu);
+	applog(LOG_WARNING, "%s", logline);
 }
 
 static void print_summary(void)
@@ -4094,6 +4163,9 @@ static void clean_up(void)
 #ifdef HAVE_OPENCL
 	clear_adl(nDevs);
 #endif
+#ifdef HAVE_LIBUSB
+        libusb_exit(NULL);
+#endif
 
 	gettimeofday(&total_tv_end, NULL);
 #ifdef HAVE_CURSES
@@ -4329,6 +4401,7 @@ void enable_curses(void) {
 	cbreak();
 	noecho();
 	curses_active = true;
+	statusy = logstart;
 	unlock_curses();
 }
 #endif
@@ -4349,6 +4422,10 @@ extern struct device_api bitforce_api;
 extern struct device_api icarus_api;
 #endif
 
+#ifdef USE_ZTEX
+extern struct device_api ztex_api;
+#endif
+
 
 static int cgminer_id_count = 0;
 
@@ -4364,7 +4441,32 @@ void enable_device(struct cgpu_info *cgpu)
 #endif
 }
 
-int main (int argc, char *argv[])
+struct _cgpu_devid_counter {
+	char name[4];
+	int lastid;
+	UT_hash_handle hh;
+};
+
+bool add_cgpu(struct cgpu_info*cgpu)
+{
+	static struct _cgpu_devid_counter *devids = NULL;
+	struct _cgpu_devid_counter *d;
+	
+	HASH_FIND_STR(devids, cgpu->api->name, d);
+	if (d)
+		cgpu->device_id = ++d->lastid;
+	else
+	{
+		d = malloc(sizeof(*d));
+		memcpy(d->name, cgpu->api->name, sizeof(d->name));
+		cgpu->device_id = d->lastid = 0;
+		HASH_ADD_STR(devids, name, d);
+	}
+	devices[total_devices++] = cgpu;
+	return true;
+}
+
+int main(int argc, char *argv[])
 {
 	struct block *block, *tmpblock;
 	struct work *work, *tmpwork;
@@ -4379,6 +4481,14 @@ int main (int argc, char *argv[])
 	if (unlikely(curl_global_init(CURL_GLOBAL_ALL)))
 		quit(1, "Failed to curl_global_init");
 
+	initial_args = malloc(sizeof(char *) * (argc + 1));
+	for  (i = 0; i < argc; i++)
+		initial_args[i] = strdup(argv[i]);
+	initial_args[argc] = NULL;
+#ifdef HAVE_LIBUSB
+        libusb_init(NULL);
+#endif
+
 	mutex_init(&hash_lock);
 	mutex_init(&qd_lock);
 #ifdef HAVE_CURSES
@@ -4531,6 +4641,10 @@ int main (int argc, char *argv[])
 	icarus_api.api_detect();
 #endif
 
+#ifdef USE_ZTEX
+	ztex_api.api_detect();
+#endif
+
 #ifdef WANT_CPUMINE
 	cpu_api.api_detect();
 #endif
@@ -4538,7 +4652,11 @@ int main (int argc, char *argv[])
 	if (devices_enabled == -1) {
 		applog(LOG_ERR, "Devices detected:");
 		for (i = 0; i < total_devices; ++i) {
-			applog(LOG_ERR, " %2d. %s%d", i, devices[i]->api->name, devices[i]->device_id);
+			struct cgpu_info *cgpu = devices[i];
+			if (cgpu->name)
+				applog(LOG_ERR, " %2d. %s %d: %s (driver: %s)", i, cgpu->api->name, cgpu->device_id, cgpu->name, cgpu->api->dname);
+			else
+				applog(LOG_ERR, " %2d. %s %d (driver: %s)", i, cgpu->api->name, cgpu->device_id, cgpu->api->dname);
 		}
 		quit(0, "%d devices listed", total_devices);
 	}

configure.ac

@@ -2,7 +2,7 @@
 ##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##
 m4_define([v_maj], [2])
 m4_define([v_min], [3])
-m4_define([v_mic], [2])
+m4_define([v_mic], [3])
 ##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##
 m4_define([v_ver], [v_maj.v_min.v_mic])
 m4_define([lt_rev], m4_eval(v_maj + v_min))
@@ -57,11 +57,14 @@ gl_INIT
 dnl Checks for header files.
 AC_HEADER_STDC
 AC_CHECK_HEADERS(syslog.h)
+AC_CHECK_HEADERS([sys/epoll.h])
 
 AC_FUNC_ALLOCA
 
 have_win32=false
 PTHREAD_FLAGS="-lpthread"
+USB_LIBS=""
+USB_FLAGS=""
 DLOPEN_FLAGS="-ldl"
 OPENCL_LIBS="-lOpenCL"
 WS2_LIBS=""
@@ -207,6 +210,16 @@ if test "x$icarus" = xyes; then
 fi
 AM_CONDITIONAL([HAS_ICARUS], [test x$icarus = xyes])
 
+ztex="no"
+
+AC_ARG_ENABLE([ztex],
+	[AC_HELP_STRING([--enable-ztex],[Compile support for Ztex (default disabled)])],
+	[ztex=$enableval]
+	)
+if test "x$ztex" = xyes; then
+	AC_DEFINE([USE_ZTEX], [1], [Defined to 1 if Ztex support is wanted])
+fi
+AM_CONDITIONAL([HAS_ZTEX], [test x$ztex = xyes])
 
 curses="auto"
 
@@ -301,6 +314,14 @@ if test "x$bitforce" != xno; then
 fi
 AM_CONDITIONAL([HAVE_LIBUDEV], [test x$libudev != xno])
 
+if test "x$ztex" != xno; then
+  AC_CHECK_LIB(usb-1.0, libusb_init, ,
+          AC_MSG_ERROR([Could not find usb library - please install libusb]))
+  AC_DEFINE([HAVE_LIBUSB], [1], [Defined to 1 if libusb is wanted])
+  USB_LIBS="-lusb-1.0"
+  USB_FLAGS=""
+fi
+
 PKG_PROG_PKG_CONFIG()
 
 PKG_CHECK_MODULES([LIBCURL], [libcurl >= 7.15.6], [AC_DEFINE([CURL_HAS_SOCKOPT], [1], [Defined if version of curl supports sockopts.])],
@@ -362,6 +383,8 @@ AC_SUBST(PDCURSES_LIBS)
 AC_SUBST(WS2_LIBS)
 AC_SUBST(MATH_LIBS)
 AC_SUBST(UDEV_LIBS)
+AC_SUBST(USB_LIBS)
+AC_SUBST(USB_FLAGS)
 
 AC_CONFIG_FILES([
 	Makefile
@@ -393,13 +416,13 @@ if test "x$opencl" != xno; then
 		echo "  OpenCL...............: FOUND. GPU mining support enabled"
 	else
 		echo "  OpenCL...............: NOT FOUND. GPU mining support DISABLED"
-		if test "x$cpumining$bitforce$icarus" = xnonono; then
+		if test "x$cpumining$bitforce$icarus$ztex" = xnononono; then
 			AC_MSG_ERROR([No mining configured in])
 		fi
 	fi
 else
 	echo "  OpenCL...............: Detection overrided. GPU mining support DISABLED"
-	if test "x$cpumining$bitforce$icarus" = xnonono; then
+	if test "x$cpumining$bitforce$icarus$ztex" = xnononono; then
 		AC_MSG_ERROR([No mining configured in])
 	fi
 fi
@@ -427,24 +450,28 @@ else
 	echo "  Icarus.FPGAs.........: Disabled"
 fi
 
+if test "x$ztex" = xyes; then
+	echo "  Ztex.FPGAs...........: Enabled"
+else
+	echo "  Ztex.FPGAs...........: Disabled"
+fi
+
 if test "x$bitforce" != xno; then
 	echo "  libudev.detection....: $libudev"
 fi
 
-echo
 if test "x$cpumining" = xyes; then
+	echo
 	echo "  CPU Mining...........: Enabled"
 	echo "  ASM.(for CPU mining).: $has_yasm"
-else
-	echo "  CPU Mining...........: Disabled"
 fi
 
 echo
 echo "Compilation............: make (or gmake)"
 echo "  CPPFLAGS.............: $CPPFLAGS"
 echo "  CFLAGS...............: $CFLAGS"
-echo "  LDFLAGS..............: $LDFLAGS $PTHREAD_FLAGS"
-echo "  LDADD................: $DLOPEN_FLAGS $LIBCURL_LIBS $JANSSON_LIBS $PTHREAD_LIBS $OPENCL_LIBS $NCURSES_LIBS $PDCURSES_LIBS $WS2_LIBS $MATH_LIBS $UDEV_LIBS"
+echo "  LDFLAGS..............: $LDFLAGS $PTHREAD_FLAGS $USB_FLAGS"
+echo "  LDADD................: $DLOPEN_FLAGS $LIBCURL_LIBS $JANSSON_LIBS $PTHREAD_LIBS $OPENCL_LIBS $NCURSES_LIBS $PDCURSES_LIBS $WS2_LIBS $MATH_LIBS $UDEV_LIBS $USB_LIBS"
 echo
 echo "Installation...........: make install (as root if needed, with 'su' or 'sudo')"
 echo "  prefix...............: $prefix"

bitforce.c → driver-bitforce.c

@@ -91,8 +91,8 @@ static void BFwrite(int fd, const void *buf, ssize_t bufLen)
 
 static bool bitforce_detect_one(const char *devpath)
 {
+	char *s;
 	char pdevbuf[0x100];
-	static int i = 0;
 
 	if (total_devices == MAX_DEVICES)
 		return false;
@@ -117,14 +117,17 @@ static bool bitforce_detect_one(const char *devpath)
 	// We have a real BitForce!
 	struct cgpu_info *bitforce;
 	bitforce = calloc(1, sizeof(*bitforce));
-	devices[total_devices++] = bitforce;
 	bitforce->api = &bitforce_api;
-	bitforce->device_id = i++;
 	bitforce->device_path = strdup(devpath);
 	bitforce->deven = DEV_ENABLED;
 	bitforce->threads = 1;
+	if (likely((!memcmp(pdevbuf, ">>>ID: ", 7)) && (s = strstr(pdevbuf + 3, ">>>"))))
+	{
+		s[0] = '\0';
+		bitforce->name = strdup(pdevbuf + 7);
+	}
 
-	return true;
+	return add_cgpu(bitforce);
 }
 
 static bool bitforce_detect_auto_udev()
@@ -358,7 +361,8 @@ static uint64_t bitforce_scanhash(struct thr_info *thr, struct work *work, uint6
 }
 
 struct device_api bitforce_api = {
-	.name = "BFL",
+	.dname = "bitforce",
+	.name = "PGA",
 	.api_detect = bitforce_detect,
 	.get_statline_before = get_bitforce_statline_before,
 	.thread_prepare = bitforce_thread_prepare,

device-cpu.c → driver-cpu.c

@@ -32,7 +32,7 @@
 #include "compat.h"
 #include "miner.h"
 #include "bench_block.h"
-#include "device-cpu.h"
+#include "driver-cpu.h"
 
 #if defined(unix)
 	#include <errno.h>
@@ -739,13 +739,13 @@ static void cpu_detect()
 	for (i = 0; i < opt_n_threads; ++i) {
 		struct cgpu_info *cgpu;
 
-		cgpu = devices[total_devices + i] = &cpus[i];
+		cgpu = &cpus[i];
 		cgpu->api = &cpu_api;
 		cgpu->deven = DEV_ENABLED;
-		cgpu->device_id = i;
 		cgpu->threads = 1;
+		cgpu->kname = algo_names[opt_algo];
+		add_cgpu(cgpu);
 	}
-	total_devices += opt_n_threads;
 }
 
 static void reinit_cpu_device(struct cgpu_info *cpu)
@@ -827,6 +827,7 @@ CPUSearch:
 }
 
 struct device_api cpu_api = {
+	.dname = "cpu",
 	.name = "CPU",
 	.api_detect = cpu_detect,
 	.reinit_device = reinit_cpu_device,

icarus.c → driver-icarus.c

@@ -46,11 +46,17 @@
   #include <windows.h>
   #include <io.h>
 #endif
+#ifdef HAVE_SYS_EPOLL_H
+  #include <sys/epoll.h>
+  #define HAVE_EPOLL
+#endif
 
 #include "elist.h"
 #include "miner.h"
 
-#define ICARUS_READ_FAULT_COUNT	(8)
+// 8 second timeout
+#define ICARUS_READ_FAULT_DECISECONDS (1)
+#define ICARUS_READ_FAULT_COUNT	(80)
 
 struct device_api icarus_api;
 
@@ -87,7 +93,7 @@ static int icarus_open(const char *devpath)
 				ISTRIP | INLCR | IGNCR | ICRNL | IXON);
 	my_termios.c_oflag &= ~OPOST;
 	my_termios.c_lflag &= ~(ECHO | ECHONL | ICANON | ISIG | IEXTEN);
-	my_termios.c_cc[VTIME] = 10; /* block 1 second */
+	my_termios.c_cc[VTIME] = ICARUS_READ_FAULT_DECISECONDS;
 	my_termios.c_cc[VMIN] = 0;
 	tcsetattr(serialfd, TCSANOW, &my_termios);
 
@@ -108,12 +114,31 @@ static int icarus_open(const char *devpath)
 #endif
 }
 
-static int icarus_gets(unsigned char *buf, size_t bufLen, int fd)
+static int icarus_gets(unsigned char *buf, size_t bufLen, int fd, volatile unsigned long *wr)
 {
 	ssize_t ret = 0;
 	int rc = 0;
+	int epollfd = -1;
+
+#ifdef HAVE_EPOLL
+	struct epoll_event ev, evr;
+	epollfd = epoll_create(1);
+	if (epollfd != -1) {
+		ev.events = EPOLLIN;
+		ev.data.fd = fd;
+		if (-1 == epoll_ctl(epollfd, EPOLL_CTL_ADD, fd, &ev)) {
+			close(epollfd);
+			epollfd = -1;
+		}
+	}
+#endif
 
 	while (bufLen) {
+#ifdef HAVE_EPOLL
+		if (epollfd != -1 && epoll_wait(epollfd, &evr, 1, ICARUS_READ_FAULT_DECISECONDS * 100) != 1)
+			ret = 0;
+		else
+#endif
 		ret = read(fd, buf, 1);
 		if (ret == 1) {
 			bufLen--;
@@ -122,13 +147,20 @@ static int icarus_gets(unsigned char *buf, size_t bufLen, int fd)
 		}
 
 		rc++;
+		if (*wr)
+			return 1;
 		if (rc == ICARUS_READ_FAULT_COUNT) {
+			if (epollfd != -1)
+				close(epollfd);
 			applog(LOG_DEBUG,
-			       "Icarus Read: No data in %d seconds", rc);
+			       "Icarus Read: No data in %d seconds", rc * ICARUS_READ_FAULT_DECISECONDS / 10);
 			return 1;
 		}
 	}
 
+	if (epollfd != -1)
+		close(epollfd);
+
 	return 0;
 }
 
@@ -172,7 +204,8 @@ static bool icarus_detect_one(const char *devpath)
 	icarus_write(fd, ob_bin, sizeof(ob_bin));
 
 	memset(nonce_bin, 0, sizeof(nonce_bin));
-	icarus_gets(nonce_bin, sizeof(nonce_bin), fd);
+	volatile unsigned long wr = 0;
+	icarus_gets(nonce_bin, sizeof(nonce_bin), fd, &wr);
 
 	icarus_close(fd);
 
@@ -194,10 +227,9 @@ static bool icarus_detect_one(const char *devpath)
 	struct cgpu_info *icarus;
 	icarus = calloc(1, sizeof(struct cgpu_info));
 	icarus->api = &icarus_api;
-	icarus->device_id = total_devices;
 	icarus->device_path = strdup(devpath);
 	icarus->threads = 1;
-	devices[total_devices++] = icarus;
+	add_cgpu(icarus);
 
 	applog(LOG_INFO, "Found Icarus at %s, mark as %d",
 	       devpath, icarus->device_id);
@@ -235,6 +267,8 @@ static bool icarus_prepare(struct thr_info *thr)
 static uint64_t icarus_scanhash(struct thr_info *thr, struct work *work,
 				__maybe_unused uint64_t max_nonce)
 {
+	volatile unsigned long *wr = &work_restart[thr->id].restart;
+
 	struct cgpu_info *icarus;
 	int fd;
 	int ret;
@@ -243,7 +277,7 @@ static uint64_t icarus_scanhash(struct thr_info *thr, struct work *work,
 	char *ob_hex, *nonce_hex;
 	uint32_t nonce;
 	uint32_t hash_count;
-	time_t t = 0;
+	struct timeval tv_start, tv_end, diff;
 
 	icarus = thr->cgpu;
 
@@ -262,6 +296,9 @@ static uint64_t icarus_scanhash(struct thr_info *thr, struct work *work,
 #ifndef WIN32
 	tcflush(fd, TCOFLUSH);
 #endif
+
+	gettimeofday(&tv_start, NULL);
+
 	ret = icarus_write(fd, ob_bin, sizeof(ob_bin));
 	if (ret) {
 		icarus_close(fd);
@@ -270,7 +307,6 @@ static uint64_t icarus_scanhash(struct thr_info *thr, struct work *work,
 
 	ob_hex = bin2hex(ob_bin, sizeof(ob_bin));
 	if (ob_hex) {
-		t = time(NULL);
 		applog(LOG_DEBUG, "Icarus %s send: %s",
 		       icarus->device_id, ob_hex);
 		free(ob_hex);
@@ -278,27 +314,35 @@ static uint64_t icarus_scanhash(struct thr_info *thr, struct work *work,
 
 	/* Icarus will return 8 bytes nonces or nothing */
 	memset(nonce_bin, 0, sizeof(nonce_bin));
-	ret = icarus_gets(nonce_bin, sizeof(nonce_bin), fd);
+	ret = icarus_gets(nonce_bin, sizeof(nonce_bin), fd, wr);
+
+	gettimeofday(&tv_end, NULL);
+	timeval_subtract(&diff, &tv_end, &tv_start);
 
 	nonce_hex = bin2hex(nonce_bin, sizeof(nonce_bin));
 	if (nonce_hex) {
-		t = time(NULL) - t;
-		applog(LOG_DEBUG, "Icarus %d return (elapse %d seconds): %s",
-		       icarus->device_id, t, nonce_hex);
+		applog(LOG_DEBUG, "Icarus %d returned (in %d.%06d seconds): %s",
+		       icarus->device_id, diff.tv_sec, diff.tv_usec, nonce_hex);
 		free(nonce_hex);
 	}
 
 	memcpy((char *)&nonce, nonce_bin, sizeof(nonce_bin));
 
+	work->blk.nonce = 0xffffffff;
+	icarus_close(fd);
+
 	if (nonce == 0 && ret) {
-		icarus_close(fd);
-                return 0xffffffff;
+		if (unlikely(diff.tv_sec > 12 || (diff.tv_sec == 11 && diff.tv_usec > 300067)))
+			return 0xffffffff;
+		// Approximately how much of the nonce Icarus scans in 1 second...
+		// 0x16a7a561 would be if it was exactly 380 MH/s
+		// 0x168b7b4b was the average over a 201-sample period based on time to find actual shares
+		return (0x168b7b4b * diff.tv_sec) + (0x17a * diff.tv_usec);
 	}
 
 #ifndef __BIG_ENDIAN__
 	nonce = swab32(nonce);
 #endif
-	work->blk.nonce = 0xffffffff;
 	submit_nonce(thr, work, nonce);
 
 	hash_count = (nonce & 0x7fffffff);
@@ -311,7 +355,8 @@ static uint64_t icarus_scanhash(struct thr_info *thr, struct work *work,
                         hash_count <<= 1;
         }
 
-	icarus_close(fd);
+	applog(LOG_DEBUG, "0x%x hashes in %d.%06d seconds", hash_count, diff.tv_sec, diff.tv_usec);
+
         return hash_count;
 }
 
@@ -333,7 +378,8 @@ static void icarus_shutdown(struct thr_info *thr)
 }
 
 struct device_api icarus_api = {
-	.name = "ICA",
+	.dname = "icarus",
+	.name = "PGA",
 	.api_detect = icarus_detect,
 	.thread_prepare = icarus_prepare,
 	.scanhash = icarus_scanhash,

device-gpu.c → driver-opencl.c

@@ -28,7 +28,7 @@
 
 #include "compat.h"
 #include "miner.h"
-#include "device-gpu.h"
+#include "driver-opencl.h"
 #include "findnonce.h"
 #include "ocl.h"
 #include "adl.h"
@@ -1107,12 +1107,13 @@ static void opencl_detect()
 	for (i = 0; i < nDevs; ++i) {
 		struct cgpu_info *cgpu;
 
-		cgpu = devices[total_devices++] = &gpus[i];
+		cgpu = &gpus[i];
 		cgpu->deven = DEV_ENABLED;
 		cgpu->api = &opencl_api;
 		cgpu->device_id = i;
 		cgpu->threads = opt_g_threads;
 		cgpu->virtual_gpu = i;
+		add_cgpu(cgpu);
 	}
 
 	if (!opt_noadl)
@@ -1179,6 +1180,7 @@ static bool opencl_thread_prepare(struct thr_info *thr)
 		return false;
 	}
 
+	strcpy(name, "");
 	applog(LOG_INFO, "Init GPU thread %i GPU %i virtual GPU %i", i, gpu, virtual_gpu);
 	clStates[i] = initCl(virtual_gpu, name, sizeof(name));
 	if (!clStates[i]) {
@@ -1208,6 +1210,26 @@ static bool opencl_thread_prepare(struct thr_info *thr)
 
 		return false;
 	}
+	if (!cgpu->name)
+		cgpu->name = strdup(name);
+	if (!cgpu->kname)
+	{
+		switch (clStates[i]->chosen_kernel) {
+		case KL_DIABLO:
+			cgpu->kname = "diablo";
+			break;
+		case KL_DIAKGCN:
+			cgpu->kname = "diakgcn";
+			break;
+		case KL_PHATK:
+			cgpu->kname = "phatk";
+			break;
+		case KL_POCLBM:
+			cgpu->kname = "poclbm";
+		default:
+			break;
+		}
+	}
 	applog(LOG_INFO, "initCl() finished. Found %s", name);
 	gettimeofday(&now, NULL);
 	get_datestamp(cgpu->init, &now);
@@ -1404,6 +1426,7 @@ static void opencl_thread_shutdown(struct thr_info *thr)
 }
 
 struct device_api opencl_api = {
+	.dname = "opencl",
 	.name = "GPU",
 	.api_detect = opencl_detect,
 	.reinit_device = reinit_opencl_device,

driver-ztex.c

@@ -0,0 +1,322 @@
+/**
+ *   ztex.c - cgminer worker for Ztex 1.15x fpga board
+ *
+ *   Copyright (c) 2012 nelisky.btc@gmail.com
+ *
+ *   This work is based upon the Java SDK provided by ztex which is
+ *   Copyright (C) 2009-2011 ZTEX GmbH.
+ *   http://www.ztex.de
+ *
+ *   This work is based upon the icarus.c worker which is
+ *   Copyright 2012 Luke Dashjr
+ *   Copyright 2012 Xiangfu <xiangfu@openmobilefree.com>
+ *
+ *   This program is free software; you can redistribute it and/or modify
+ *   it under the terms of the GNU General Public License version 2 as
+ *   published by the Free Software Foundation.
+ *
+ *   This program is distributed in the hope that it will be useful, but
+ *   WITHOUT ANY WARRANTY; without even the implied warranty of
+ *   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ *   General Public License for more details.
+ *
+ *   You should have received a copy of the GNU General Public License
+ *   along with this program; if not, see http://www.gnu.org/licenses/.
+**/
+#include <unistd.h>
+#include <sha2.h>
+#include "miner.h"
+#include "libztex.h"
+
+#define GOLDEN_BACKLOG 5
+
+struct device_api ztex_api, ztex_hotplug_api;
+
+// Forward declarations
+static void ztex_disable(struct thr_info* thr);
+static bool ztex_prepare(struct thr_info *thr);
+
+static void ztex_detect(void)
+{
+	int cnt;
+	int i;
+	struct libztex_dev_list **ztex_devices;
+
+	cnt = libztex_scanDevices(&ztex_devices);
+	applog(LOG_WARNING, "Found %d ztex board(s)", cnt);
+
+	for (i = 0; i < cnt; i++) {
+		if (total_devices == MAX_DEVICES)
+			break;
+		struct cgpu_info *ztex;
+		ztex = calloc(1, sizeof(struct cgpu_info));
+		ztex->api = &ztex_api;
+		ztex->device_ztex = ztex_devices[i]->dev;
+		ztex->threads = 1;
+		add_cgpu(ztex);
+
+		applog(LOG_WARNING,"%s: Found Ztex, mark as %d", ztex->device_ztex->repr, ztex->device_id);
+	}
+
+	if (cnt > 0)
+		libztex_freeDevList(ztex_devices);
+}
+
+static bool ztex_updateFreq(struct libztex_device* ztex)
+{
+	int i, maxM, bestM;
+	double bestR, r;
+
+	for (i = 0; i < ztex->freqMaxM; i++)
+		if (ztex->maxErrorRate[i + 1] * i < ztex->maxErrorRate[i] * (i + 20))
+			ztex->maxErrorRate[i + 1] = ztex->maxErrorRate[i] * (1.0 + 20.0 / i);
+
+	maxM = 0;
+	while (maxM < ztex->freqMDefault && ztex->maxErrorRate[maxM + 1] < LIBZTEX_MAXMAXERRORRATE)
+		maxM++;
+	while (maxM < ztex->freqMaxM && ztex->errorWeight[maxM] > 150 && ztex->maxErrorRate[maxM + 1] < LIBZTEX_MAXMAXERRORRATE)
+		maxM++;
+
+	bestM = 0;
+	bestR = 0;
+	for (i = 0; i <= maxM; i++) {
+		r = (i + 1 + (i == ztex->freqM? LIBZTEX_ERRORHYSTERESIS: 0)) * (1 - ztex->maxErrorRate[i]);
+		if (r > bestR) {
+			bestM = i;
+			bestR = r;
+		}
+	}
+
+	if (bestM != ztex->freqM) 
+		libztex_setFreq(ztex, bestM);
+
+	maxM = ztex->freqMDefault;
+	while (maxM < ztex->freqMaxM && ztex->errorWeight[maxM + 1] > 100)
+		maxM++;
+	if ((bestM < (1.0 - LIBZTEX_OVERHEATTHRESHOLD) * maxM) && bestM < maxM - 1) {
+		libztex_resetFpga(ztex);
+		applog(LOG_ERR, "%s: frequency drop of %.1f%% detect. This may be caused by overheating. FPGA is shut down to prevent damage.",
+		       ztex->repr, (1.0 - 1.0 * bestM / maxM) * 100);
+		return false;
+	}
+	return true;
+}
+
+
+static bool ztex_checkNonce(struct libztex_device *ztex,
+                            struct work *work,
+                            struct libztex_hash_data *hdata)
+{
+	uint32_t *data32 = (uint32_t *)(work->data);
+	unsigned char swap[128];
+	uint32_t *swap32 = (uint32_t *)swap;
+	unsigned char hash1[32];
+	unsigned char hash2[32];
+	uint32_t *hash2_32 = (uint32_t *)hash2;
+	int i;
+
+#if defined(__BIGENDIAN__) || defined(MIPSEB)
+	hdata->nonce = swab32(hdata->nonce);
+	hdata->hash7 = swab32(hdata->hash7);
+#endif
+
+	work->data[64 + 12 + 0] = (hdata->nonce >> 0) & 0xff;
+	work->data[64 + 12 + 1] = (hdata->nonce >> 8) & 0xff;
+	work->data[64 + 12 + 2] = (hdata->nonce >> 16) & 0xff;
+	work->data[64 + 12 + 3] = (hdata->nonce >> 24) & 0xff;
+
+	for (i = 0; i < 80 / 4; i++)
+		swap32[i] = swab32(data32[i]);
+	
+	sha2(swap, 80, hash1, false);
+	sha2(hash1, 32, hash2, false);
+#if defined(__BIGENDIAN__) || defined(MIPSEB)
+	if (hash2_32[7] != ((hdata->hash7 + 0x5be0cd19) & 0xFFFFFFFF)) {
+#else
+	if (swab32(hash2_32[7]) != ((hdata->hash7 + 0x5be0cd19) & 0xFFFFFFFF)) {
+#endif
+		ztex->errorCount[ztex->freqM] += 1.0 / ztex->numNonces;
+		applog(LOG_DEBUG, "%s: checkNonce failed for %0.8X", ztex->repr, hdata->nonce);
+		return false;
+	}
+	return true;
+}
+
+static uint64_t ztex_scanhash(struct thr_info *thr, struct work *work,
+                              __maybe_unused uint64_t max_nonce)
+{
+	struct libztex_device *ztex;
+	unsigned char sendbuf[44];
+	int i, j;
+	uint32_t backlog[GOLDEN_BACKLOG];
+	int backlog_p = 0;
+	uint32_t lastnonce[GOLDEN_BACKLOG], nonce, noncecnt = 0;
+	bool overflow, found, rv;
+	struct libztex_hash_data hdata[GOLDEN_BACKLOG];
+
+	ztex = thr->cgpu->device_ztex;
+
+	memcpy(sendbuf, work->data + 64, 12);
+	memcpy(sendbuf + 12, work->midstate, 32);
+	memset(backlog, 0, sizeof(backlog));
+	i = libztex_sendHashData(ztex, sendbuf);
+	if (i < 0) {
+		// Something wrong happened in send
+		applog(LOG_ERR, "%s: Failed to send hash data with err %d, retrying", ztex->repr, i);
+		usleep(500000);
+		i = libztex_sendHashData(ztex, sendbuf);
+		if (i < 0) {
+			// And there's nothing we can do about it
+			ztex_disable(thr);
+			applog(LOG_ERR, "%s: Failed to send hash data with err %d, giving up", ztex->repr, i);
+			return 0;
+		}
+	}
+	
+	applog(LOG_DEBUG, "sent hashdata");
+
+	for (i = 0; i < ztex->numNonces; i++)
+		lastnonce[i] = 0;
+
+	overflow = false;
+
+	while (!(overflow || work_restart[thr->id].restart)) {
+		usleep(250000);
+		if (work_restart[thr->id].restart) {
+			applog(LOG_DEBUG, "%s: New work detected", ztex->repr);
+			break;
+		}
+		i = libztex_readHashData(ztex, &hdata[0]);
+		if (i < 0) {
+			// Something wrong happened in read
+			applog(LOG_ERR, "%s: Failed to read hash data with err %d, retrying", ztex->repr, i);
+			usleep(500000);
+			i = libztex_readHashData(ztex, &hdata[0]);
+			if (i < 0) {
+				// And there's nothing we can do about it
+				ztex_disable(thr);
+				applog(LOG_ERR, "%s: Failed to read hash data with err %d, giving up", ztex->repr, i);
+				return 0;
+			}
+		}
+
+		if (work_restart[thr->id].restart) {
+			applog(LOG_DEBUG, "%s: New work detected", ztex->repr);
+			break;
+		}
+
+		ztex->errorCount[ztex->freqM] *= 0.995;
+		ztex->errorWeight[ztex->freqM] = ztex->errorWeight[ztex->freqM] * 0.995 + 1.0;
+ 
+		for (i = 0; i < ztex->numNonces; i++) {
+			nonce = hdata[i].nonce;
+#if defined(__BIGENDIAN__) || defined(MIPSEB)
+			nonce = swab32(nonce);
+#endif
+			if (nonce > noncecnt)
+				noncecnt = nonce;
+			if (((nonce & 0x7fffffff) >> 4) < ((lastnonce[i] & 0x7fffffff) >> 4)) {
+				applog(LOG_DEBUG, "%s: overflow nonce=%0.8x lastnonce=%0.8x", ztex->repr, nonce, lastnonce[i]);
+				overflow = true;
+			} else
+				lastnonce[i] = nonce;
+#if !(defined(__BIGENDIAN__) || defined(MIPSEB))
+			nonce = swab32(nonce);
+#endif
+			if (!ztex_checkNonce(ztex, work, &hdata[i])) {
+				thr->cgpu->hw_errors++;
+				continue;
+			}
+			nonce = hdata[i].goldenNonce;
+			if (nonce > 0) {
+				found = false;
+				for (j = 0; j < GOLDEN_BACKLOG; j++) {
+					if (backlog[j] == nonce) {
+						found = true;
+						break;
+					}
+				}
+				if (!found) {
+					applog(LOG_DEBUG, "%s: Share found", ztex->repr);
+					backlog[backlog_p++] = nonce;
+					if (backlog_p >= GOLDEN_BACKLOG)
+						backlog_p = 0;
+#if defined(__BIGENDIAN__) || defined(MIPSEB)
+					nonce = swab32(nonce);
+#endif
+					work->blk.nonce = 0xffffffff;
+					rv = submit_nonce(thr, work, nonce);
+					applog(LOG_DEBUG, "%s: submitted %0.8x %d", ztex->repr, nonce, rv);
+				}
+			}
+
+		}
+
+	}
+
+	ztex->errorRate[ztex->freqM] = ztex->errorCount[ztex->freqM] /	ztex->errorWeight[ztex->freqM] * (ztex->errorWeight[ztex->freqM] < 100? ztex->errorWeight[ztex->freqM] * 0.01: 1.0);
+	if (ztex->errorRate[ztex->freqM] > ztex->maxErrorRate[ztex->freqM])
+		ztex->maxErrorRate[ztex->freqM] = ztex->errorRate[ztex->freqM];
+
+	if (!ztex_updateFreq(ztex))
+		// Something really serious happened, so mark this thread as dead!
+		return 0;
+
+	applog(LOG_DEBUG, "%s: exit %1.8X", ztex->repr, noncecnt);
+
+	work->blk.nonce = 0xffffffff;
+
+	return noncecnt > 0? noncecnt: 1;
+}
+
+static void ztex_statline_before(char *buf, struct cgpu_info *cgpu)
+{
+	if (cgpu->deven == DEV_ENABLED) {
+		tailsprintf(buf, "%s | ", cgpu->device_ztex->snString);
+		tailsprintf(buf, "%0.2fMhz | ", cgpu->device_ztex->freqM1 * (cgpu->device_ztex->freqM + 1));
+	}
+}
+
+static bool ztex_prepare(struct thr_info *thr)
+{
+	struct timeval now;
+	struct cgpu_info *ztex = thr->cgpu;
+
+	gettimeofday(&now, NULL);
+	get_datestamp(ztex->init, &now);
+
+	if (libztex_configureFpga(ztex->device_ztex) != 0)
+		return false;
+
+	ztex->device_ztex->freqM = -1;
+	ztex_updateFreq(ztex->device_ztex);
+
+	applog(LOG_DEBUG, "%s: prepare", ztex->device_ztex->repr);
+	return true;
+}
+
+static void ztex_shutdown(struct thr_info *thr)
+{
+	if (thr->cgpu->device_ztex != NULL) {
+		applog(LOG_DEBUG, "%s: shutdown", thr->cgpu->device_ztex->repr);
+		libztex_destroy_device(thr->cgpu->device_ztex);
+		thr->cgpu->device_ztex = NULL;
+	}
+}
+
+static void ztex_disable(struct thr_info *thr)
+{
+	applog(LOG_ERR, "%s: Disabling!", thr->cgpu->device_ztex->repr);
+	devices[thr->cgpu->device_id]->deven = DEV_DISABLED;
+	ztex_shutdown(thr);
+}
+
+struct device_api ztex_api = {
+	.dname = "ztex",
+	.name = "PGA",
+	.api_detect = ztex_detect,
+	.get_statline_before = ztex_statline_before,
+	.thread_prepare = ztex_prepare,
+	.scanhash = ztex_scanhash,
+	.thread_shutdown = ztex_shutdown,
+};

libztex.c

@@ -0,0 +1,477 @@
+/**
+ *   libztex.c - Ztex 1.15x fpga board support library
+ *
+ *   Copyright (c) 2012 nelisky.btc@gmail.com
+ *
+ *   This work is based upon the Java SDK provided by ztex which is
+ *   Copyright (C) 2009-2011 ZTEX GmbH.
+ *   http://www.ztex.de
+ *
+ *   This program is free software; you can redistribute it and/or modify
+ *   it under the terms of the GNU General Public License version 2 as
+ *   published by the Free Software Foundation.
+ *
+ *   This program is distributed in the hope that it will be useful, but
+ *   WITHOUT ANY WARRANTY; without even the implied warranty of
+ *   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ *   General Public License for more details.
+ *
+ *   You should have received a copy of the GNU General Public License
+ *   along with this program; if not, see http://www.gnu.org/licenses/.
+**/
+
+#include <stdio.h>
+#include <unistd.h>
+#include "miner.h"
+#include "libztex.h"
+
+#define BUFSIZE 256
+
+//* Capability index for EEPROM support.
+#define CAPABILITY_EEPROM 0,0
+//* Capability index for FPGA configuration support. 
+#define CAPABILITY_FPGA 0,1
+//* Capability index for FLASH memory support.
+#define CAPABILITY_FLASH 0,2
+//* Capability index for DEBUG helper support.
+#define CAPABILITY_DEBUG 0,3
+//* Capability index for AVR XMEGA support.
+#define CAPABILITY_XMEGA 0,4
+//* Capability index for AVR XMEGA support.
+#define CAPABILITY_HS_FPGA 0,5
+//* Capability index for AVR XMEGA support.
+#define CAPABILITY_MAC_EEPROM 0,6
+
+
+
+static bool libztex_checkDevice(struct libusb_device *dev)
+{
+	struct libusb_device_descriptor desc;
+	int err;
+
+	err = libusb_get_device_descriptor(dev, &desc);
+	if (unlikely(err != 0)) {
+		applog(LOG_ERR, "Ztex check device: Failed to open read descriptor with error %d", err);
+		return false;
+	}
+	if (!(desc.idVendor == LIBZTEX_IDVENDOR && desc.idProduct == LIBZTEX_IDPRODUCT)) {
+		applog(LOG_DEBUG, "Not a ZTEX device %0.4x:%0.4x", desc.idVendor, desc.idProduct);
+		return false;
+	}
+	return true;
+}
+
+static bool libztex_checkCapability(struct libztex_device *ztex, int i, int j)
+{
+	if (!((i >= 0) && (i <= 5) && (j >= 0) && (j < 8) &&
+	     (((ztex->interfaceCapabilities[i] & 255) & (1 << j)) != 0))) {
+		applog(LOG_ERR, "%s: capability missing: %d %d", ztex->repr, i, i);
+		return false;
+	}
+	return true;
+}
+
+static int libztex_detectBitstreamBitOrder(const unsigned char *buf, int size)
+{
+	int i;
+
+	for (i = 0; i < size - 4; i++) {
+		if (((buf[i] & 255) == 0xaa) && ((buf[i + 1] & 255) == 0x99) && ((buf[i + 2] & 255) == 0x55) && ((buf[i + 3] & 255) == 0x66))
+			return 1;
+		if (((buf[i] & 255) == 0x55) && ((buf[i + 1] & 255) == 0x99) && ((buf[i + 2] & 255) == 0xaa) && ((buf[i + 3] & 255) == 0x66))
+			return 0;
+	} 
+	applog(LOG_WARNING, "Unable to determine bitstream bit order: no signature found");
+	return 0;
+}
+
+static void libztex_swapBits(unsigned char *buf, int size)
+{
+	unsigned char c;
+	int i;
+
+	for (i = 0; i < size; i++) {
+		c = buf[i];
+		buf[i] = ((c & 128) >> 7) |
+		         ((c & 64) >> 5) |
+		         ((c & 32) >> 3) |
+		         ((c & 16) >> 1) |
+		         ((c & 8) << 1) |
+		         ((c & 4) << 3) |
+		         ((c & 2) << 5) |
+		         ((c & 1) << 7);
+	}
+}
+
+static int libztex_getFpgaState(struct libztex_device *ztex, struct libztex_fpgastate *state)
+{
+	unsigned char buf[9];
+	int cnt;
+
+	if (!libztex_checkCapability(ztex, CAPABILITY_FPGA))
+		return -1;
+	cnt = libusb_control_transfer(ztex->hndl, 0xc0, 0x30, 0, 0, buf, 9, 1000);
+	if (unlikely(cnt < 0)) {
+		applog(LOG_ERR, "%s: Failed getFpgaState with err %d", ztex->repr, cnt);
+		return cnt;
+	}
+	state->fpgaConfigured = (buf[0] == 0);
+	state->fpgaChecksum = buf[1] & 0xff;
+	state->fpgaBytes = ((buf[5] & 0xff) << 24) | ((buf[4] & 0xff) << 16) | ((buf[3] & 0xff) << 8) | (buf[2] & 0xff);
+	state->fpgaInitB = buf[6] & 0xff;
+	state->fpgaFlashResult = buf[7];
+	state->fpgaFlashBitSwap = (buf[8] != 0);
+	return 0;
+}
+
+static int libztex_configureFpgaLS(struct libztex_device *ztex, const char* firmware, bool force, char bs)
+{
+	struct libztex_fpgastate state;
+	const int transactionBytes = 2048;
+	unsigned char buf[transactionBytes], cs;
+	int tries, cnt, buf_p, i;
+	ssize_t pos = 0;
+	FILE *fp;
+
+	if (!libztex_checkCapability(ztex, CAPABILITY_FPGA))
+		return -1;
+
+	libztex_getFpgaState(ztex, &state);
+	if (!force && state.fpgaConfigured) {
+		applog(LOG_DEBUG, "Bitstream already configured");
+		return 1;
+	}
+
+	for (tries = 10; tries > 0; tries--) {
+		fp = fopen(firmware, "rb");
+		if (!fp) {
+			applog(LOG_ERR, "%s: failed to read firmware '%s'", ztex->repr, firmware);
+			return -2;
+		}
+
+		cs = 0;
+		while (pos < transactionBytes && !feof(fp)) {
+			buf[pos] = getc(fp);
+			cs += buf[pos++];
+		}
+
+		if (feof(fp))
+			pos--;
+
+		if (bs != 0 && bs != 1)
+			bs = libztex_detectBitstreamBitOrder(buf, transactionBytes < pos? transactionBytes: pos);
+
+		//* Reset fpga
+		cnt = libztex_resetFpga(ztex);
+		if (unlikely(cnt < 0)) {
+			applog(LOG_ERR, "%s: Failed reset fpga with err %d", ztex->repr, cnt);
+			continue;
+		}
+
+		if (bs == 1)
+			libztex_swapBits(buf, pos);
+	 
+		buf_p = pos;
+		while (1) {
+			i = 0;
+			while (i < buf_p) {
+				cnt = libusb_control_transfer(ztex->hndl, 0x40, 0x32, 0, 0, &buf[i], buf_p - i, 5000);
+				if (unlikely(cnt < 0)) {
+					applog(LOG_ERR, "%s: Failed send fpga data with err %d", ztex->repr, cnt);
+					break;
+				}
+				i += cnt;
+			}
+			if (i < buf_p || buf_p < transactionBytes)
+				break;
+			buf_p = 0;
+			while (buf_p < transactionBytes && !feof(fp)) {
+				buf[buf_p] = getc(fp);
+				cs += buf[buf_p++];
+			}
+			if (feof(fp))
+				buf_p--;
+			pos += buf_p;
+			if (buf_p == 0)
+				break;
+			if (bs == 1)
+				libztex_swapBits(buf, buf_p);
+		}
+		if (cnt >= 0)
+			tries = 0;
+
+		fclose(fp);
+	}
+	libztex_getFpgaState(ztex, &state);
+	if (!state.fpgaConfigured) {
+		applog(LOG_ERR, "%s: FPGA configuration failed: DONE pin does not go high", ztex->repr);
+		return 3;
+	}
+	usleep(200000);
+	applog(LOG_INFO, "%s: FPGA configuration done", ztex->repr);
+	return 0;
+}
+
+int libztex_configureFpga(struct libztex_device *ztex)
+{
+	char buf[256] = "bitstreams/";
+
+	memset(&buf[11], 0, 245);
+	strcpy(&buf[11], ztex->bitFileName);
+	strcpy(&buf[strlen(buf)], ".bit");
+
+	return libztex_configureFpgaLS(ztex, buf, true, 2);
+}
+
+int libztex_setFreq(struct libztex_device *ztex, uint16_t freq)
+{
+	int cnt;
+
+	if (freq > ztex->freqMaxM)
+		freq = ztex->freqMaxM;
+
+	cnt = libusb_control_transfer(ztex->hndl, 0x40, 0x83, freq, 0, NULL, 0, 500);
+	if (unlikely(cnt < 0)) {
+		applog(LOG_ERR, "Ztex check device: Failed to set frequency with err %d", cnt);
+		return cnt;
+	}
+	ztex->freqM = freq;
+	applog(LOG_WARNING, "%s: Frequency change to %0.2f Mhz", ztex->repr, ztex->freqM1 * (ztex->freqM + 1));
+
+	return 0;
+}
+
+int libztex_resetFpga(struct libztex_device *ztex)
+{
+	return libusb_control_transfer(ztex->hndl, 0x40, 0x31, 0, 0, NULL, 0, 1000);
+}
+
+int libztex_prepare_device(struct libusb_device *dev, struct libztex_device** ztex)
+{
+	struct libztex_device *newdev;
+	unsigned char buf[64];
+	int cnt, err;
+
+	newdev = malloc(sizeof(struct libztex_device));
+	newdev->bitFileName = NULL;
+	newdev->valid = false;
+	newdev->hndl = NULL;
+	*ztex = newdev;
+
+	err = libusb_get_device_descriptor(dev, &newdev->descriptor);
+	if (unlikely(err != 0)) {
+		applog(LOG_ERR, "Ztex check device: Failed to open read descriptor with error %d", err);
+		return err;
+	}
+
+	// Check vendorId and productId
+	if (!(newdev->descriptor.idVendor == LIBZTEX_IDVENDOR &&
+				newdev->descriptor.idProduct == LIBZTEX_IDPRODUCT)) {
+		applog(LOG_ERR, "Not a ztex device? %0.4X, %0.4X", newdev->descriptor.idVendor, newdev->descriptor.idProduct);
+		return 1;
+	}
+
+	err = libusb_open(dev, &newdev->hndl);
+	if (unlikely(err != 0)) {
+		applog(LOG_ERR, "Ztex check device: Failed to open handle with error %d", err);
+		return err;
+	}
+
+	cnt = libusb_get_string_descriptor_ascii (newdev->hndl, newdev->descriptor.iSerialNumber, newdev->snString,
+	                                          LIBZTEX_SNSTRING_LEN + 1);
+	if (unlikely(cnt < 0)) {
+		applog(LOG_ERR, "Ztex check device: Failed to read device snString with err %d", cnt);
+		return cnt;
+	}
+
+	cnt = libusb_control_transfer(newdev->hndl, 0xc0, 0x22, 0, 0, buf, 40, 500);
+	if (unlikely(cnt < 0)) {
+		applog(LOG_ERR, "Ztex check device: Failed to read ztex descriptor with err %d", cnt);
+		return cnt;
+	}
+	
+	if ( buf[0] != 40 || buf[1] != 1 || buf[2] != 'Z' || buf[3] != 'T' || buf[4] != 'E' || buf[5] != 'X' ) {
+		applog(LOG_ERR, "Ztex check device: Error reading ztex descriptor");
+		return 2;
+	}
+
+	newdev->productId[0] = buf[6];
+	newdev->productId[1] = buf[7];
+	newdev->productId[2] = buf[8];
+	newdev->productId[3] = buf[9];
+	newdev->fwVersion = buf[10];
+	newdev->interfaceVersion = buf[11];
+	newdev->interfaceCapabilities[0] = buf[12];
+	newdev->interfaceCapabilities[1] = buf[13];
+	newdev->interfaceCapabilities[2] = buf[14];
+	newdev->interfaceCapabilities[3] = buf[15];
+	newdev->interfaceCapabilities[4] = buf[16];
+	newdev->interfaceCapabilities[5] = buf[17];
+	newdev->moduleReserved[0] = buf[18];
+	newdev->moduleReserved[1] = buf[19];
+	newdev->moduleReserved[2] = buf[20];
+	newdev->moduleReserved[3] = buf[21];
+	newdev->moduleReserved[4] = buf[22];
+	newdev->moduleReserved[5] = buf[23];
+	newdev->moduleReserved[6] = buf[24];
+	newdev->moduleReserved[7] = buf[25];
+	newdev->moduleReserved[8] = buf[26];
+	newdev->moduleReserved[9] = buf[27];
+	newdev->moduleReserved[10] = buf[28];
+	newdev->moduleReserved[11] = buf[29];
+
+
+	cnt = libusb_control_transfer(newdev->hndl, 0xc0, 0x82, 0, 0, buf, 64, 500);
+	if (unlikely(cnt < 0)) {
+		applog(LOG_ERR, "Ztex check device: Failed to read ztex descriptor with err %d", cnt);
+		return cnt;
+	}
+
+	if (unlikely(buf[0] != 4)) {
+		if (unlikely(buf[0] != 2)) {
+			applog(LOG_ERR, "Invalid BTCMiner descriptor version. Firmware must be updated (%d).", buf[0]);
+			return 3;
+		}
+		applog(LOG_WARNING, "Firmware out of date");
+	}
+
+	newdev->numNonces = buf[1] + 1;
+	newdev->offsNonces = ((buf[2] & 255) | ((buf[3] & 255) << 8)) - 10000;
+	newdev->freqM1 = ((buf[4] & 255) | ((buf[5] & 255) << 8) ) * 0.01;
+	newdev->freqMaxM = (buf[7] & 255);
+	newdev->freqM = (buf[6] & 255);
+	newdev->freqMDefault = newdev->freqM;
+
+	for (cnt=0; cnt < 255; cnt++) {
+		newdev->errorCount[cnt] = 0;
+		newdev->errorWeight[cnt] = 0;
+		newdev->errorRate[cnt] = 0;
+		newdev->maxErrorRate[cnt] = 0;
+	}
+
+	cnt = strlen((char *)&buf[buf[0] == 4? 10: 8]);
+	newdev->bitFileName = malloc(sizeof(char) * (cnt + 1));
+	memcpy(newdev->bitFileName, &buf[buf[0] == 4? 10: 8], cnt + 1);
+
+	newdev->usbbus = libusb_get_bus_number(dev);
+	newdev->usbaddress = libusb_get_device_address(dev);
+	sprintf(newdev->repr, "ZTEX %.3d:%.3d-%s", newdev->usbbus, newdev->usbaddress, newdev->snString);
+	newdev->valid = true;
+	return 0;
+}
+
+void libztex_destroy_device(struct libztex_device* ztex)
+{
+	if (ztex->hndl != NULL) {
+		libusb_close(ztex->hndl);
+		ztex->hndl = NULL;
+	}
+	if (ztex->bitFileName != NULL) {
+		free(ztex->bitFileName);
+		ztex->bitFileName = NULL;
+	}
+	free(ztex);
+}
+
+int libztex_scanDevices(struct libztex_dev_list*** devs_p)
+{
+	int usbdevices[LIBZTEX_MAX_DESCRIPTORS];
+	struct libztex_dev_list **devs;
+	struct libztex_device *ztex;
+	int found = 0, pos = 0, err;
+	libusb_device **list;
+	ssize_t cnt, i = 0;
+
+	cnt = libusb_get_device_list(NULL, &list);
+	if (unlikely(cnt < 0)) {
+		applog(LOG_ERR, "Ztex scan devices: Failed to list usb devices with err %d", cnt);
+		return 0;
+	}
+
+	for (i = 0; i < cnt; i++) {
+		if (libztex_checkDevice(list[i])) {
+			// Got one!
+			usbdevices[found] = i;
+			found++;
+		}
+	}
+
+	devs = malloc(sizeof(struct libztex_dev_list *) * found);
+	if (devs == NULL) {
+		applog(LOG_ERR, "Ztex scan devices: Failed to allocate memory");
+		return 0;
+	}
+
+	for (i = 0; i < found; i++) {
+		err = libztex_prepare_device(list[usbdevices[i]], &ztex);
+		if (unlikely(err != 0))
+			applog(LOG_ERR, "prepare device: %d", err);
+		// check if valid
+		if (!ztex->valid) {
+			libztex_destroy_device(ztex);
+			continue;
+		}
+		devs[pos] = malloc(sizeof(struct libztex_dev_list));
+		devs[pos]->dev = ztex;
+		devs[pos]->next = NULL;
+		if (pos > 0)
+			devs[pos - 1]->next = devs[pos];
+		pos++;
+	}
+
+	libusb_free_device_list(list, 1);
+	*devs_p = devs;
+	return pos;
+}
+
+int libztex_sendHashData(struct libztex_device *ztex, unsigned char *sendbuf)
+{
+	int cnt;
+
+	if (ztex == NULL || ztex->hndl == NULL)
+		return 0;
+	cnt = libusb_control_transfer(ztex->hndl, 0x40, 0x80, 0, 0, sendbuf, 44, 1000);
+	if (unlikely(cnt < 0))
+		applog(LOG_ERR, "%s: Failed sendHashData with err %d", ztex->repr, cnt);
+	
+	return cnt;
+}
+
+int libztex_readHashData(struct libztex_device *ztex, struct libztex_hash_data nonces[])
+{
+	// length of buf must be 8 * (numNonces + 1)
+	unsigned char rbuf[12 * 8];
+	int cnt, i;
+
+	if (ztex->hndl == NULL)
+		return 0;
+	
+	cnt = libusb_control_transfer(ztex->hndl, 0xc0, 0x81, 0, 0, rbuf, 12 * ztex->numNonces, 1000);
+	if (unlikely(cnt < 0)) {
+		applog(LOG_ERR, "%s: Failed readHashData with err %d", ztex->repr, cnt);
+		return cnt;
+	}
+
+	for (i = 0; i < ztex->numNonces; i++) {
+		memcpy((char*)&nonces[i].goldenNonce, &rbuf[i * 12], 4);
+		nonces[i].goldenNonce -= ztex->offsNonces;
+		memcpy((char*)&nonces[i].nonce, &rbuf[(i * 12) + 4], 4);
+		nonces[i].nonce -= ztex->offsNonces;
+		memcpy((char*)&nonces[i].hash7, &rbuf[(i * 12) + 8], 4);
+	}
+	
+	return cnt;
+}
+
+void libztex_freeDevList(struct libztex_dev_list **devs)
+{
+	bool done = false;
+	ssize_t cnt = 0;
+
+	while (!done) {
+		if (devs[cnt]->next == NULL)
+			done = true;
+		free(devs[cnt++]);
+	}
+	free(devs);
+}

libztex.h

@@ -0,0 +1,95 @@
+/**
+ *   libztex.h - headers for Ztex 1.15x fpga board support library
+ *
+ *   Copyright (c) 2012 nelisky.btc@gmail.com
+ *
+ *   This work is based upon the Java SDK provided by ztex which is
+ *   Copyright (C) 2009-2011 ZTEX GmbH.
+ *   http://www.ztex.de
+ *
+ *   This program is free software; you can redistribute it and/or modify
+ *   it under the terms of the GNU General Public License version 2 as
+ *   published by the Free Software Foundation.
+ *
+ *   This program is distributed in the hope that it will be useful, but
+ *   WITHOUT ANY WARRANTY; without even the implied warranty of
+ *   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ *   General Public License for more details.
+ *
+ *   You should have received a copy of the GNU General Public License
+ *   along with this program; if not, see http://www.gnu.org/licenses/.
+**/
+#ifndef __LIBZTEX_H__
+#define __LIBZTEX_H__
+
+#include <libusb-1.0/libusb.h>
+
+#define LIBZTEX_MAX_DESCRIPTORS 512
+#define LIBZTEX_SNSTRING_LEN 10
+
+#define LIBZTEX_IDVENDOR 0x221A
+#define LIBZTEX_IDPRODUCT 0x0100
+
+#define LIBZTEX_MAXMAXERRORRATE 0.05
+#define LIBZTEX_ERRORHYSTERESIS 0.1
+#define LIBZTEX_OVERHEATTHRESHOLD 0.5
+
+struct libztex_fpgastate {
+	bool fpgaConfigured;
+	unsigned char fpgaChecksum;
+	uint16_t fpgaBytes;
+	unsigned char fpgaInitB;
+	unsigned char fpgaFlashResult;
+	bool fpgaFlashBitSwap;
+};
+
+struct libztex_device {
+	bool valid;
+	struct libusb_device_descriptor descriptor;
+	libusb_device_handle *hndl; 
+	unsigned char usbbus;
+	unsigned char usbaddress;
+	unsigned char snString[LIBZTEX_SNSTRING_LEN+1];
+	unsigned char productId[4];
+	unsigned char fwVersion;
+	unsigned char interfaceVersion;
+	unsigned char interfaceCapabilities[6];
+	unsigned char moduleReserved[12];
+	uint8_t numNonces;
+	uint16_t offsNonces;
+	double freqM1;	
+	uint8_t freqM;
+	uint8_t freqMaxM;
+	uint8_t freqMDefault;
+	char* bitFileName;
+
+	double errorCount[256];
+	double errorWeight[256];
+	double errorRate[256];
+	double maxErrorRate[256];
+
+	char repr[64];
+};
+
+struct libztex_dev_list { 
+	struct libztex_device *dev;
+	struct libztex_dev_list *next;
+};
+
+struct libztex_hash_data {
+	uint32_t goldenNonce;
+	uint32_t nonce;
+	uint32_t hash7;
+};
+
+extern int libztex_scanDevices (struct libztex_dev_list ***devs);
+extern void libztex_freeDevList (struct libztex_dev_list **devs);
+extern int libztex_prepare_device (struct libusb_device *dev, struct libztex_device** ztex);
+extern void libztex_destroy_device (struct libztex_device* ztex);
+extern int libztex_configureFpga (struct libztex_device *dev);
+extern int libztex_setFreq (struct libztex_device *ztex, uint16_t freq);
+extern int libztex_sendHashData (struct libztex_device *ztex, unsigned char *sendbuf);
+extern int libztex_readHashData (struct libztex_device *ztex, struct libztex_hash_data nonces[]);
+extern int libztex_resetFpga (struct libztex_device *ztex);
+
+#endif /* __LIBZTEX_H__ */

miner.h

@@ -61,6 +61,14 @@ void *alloca (size_t);
  #include "ADL_SDK/adl_sdk.h"
 #endif
 
+#ifdef HAVE_LIBUSB
+  #include <libusb-1.0/libusb.h>
+#endif
+
+#ifdef USE_ZTEX
+  #include "libztex.h"
+#endif
+
 #if !defined(WIN32) && ((__GNUC__ > 4) || (__GNUC__ == 4 && __GNUC_MINOR__ >= 3))
 #define bswap_16 __builtin_bswap16
 #define bswap_32 __builtin_bswap32
@@ -187,6 +195,7 @@ struct thr_info;
 struct work;
 
 struct device_api {
+	char*dname;
 	char*name;
 
 	// API-global functions
@@ -247,9 +256,15 @@ struct cgpu_info {
 	int cgminer_id;
 	struct device_api *api;
 	int device_id;
+	char *name;
 	char *device_path;
 	FILE *device_file;
-	int device_fd;
+	union {
+#ifdef USE_ZTEX
+		struct libztex_device *device_ztex;
+#endif
+		int device_fd;
+	};
 
 	enum dev_enable deven;
 	int accepted;
@@ -270,6 +285,7 @@ struct cgpu_info {
 	int virtual_gpu;
 	int intensity;
 	bool dynamic;
+	char *kname;
 #ifdef HAVE_OPENCL
 	cl_uint vwidth;
 	size_t work_size;
@@ -308,6 +324,8 @@ struct cgpu_info {
 	int dev_thermal_cutoff_count;
 };
 
+extern bool add_cgpu(struct cgpu_info*);
+
 struct thread_q {
 	struct list_head	q;
 
@@ -645,6 +663,7 @@ extern int curses_int(const char *query);
 extern char *curses_input(const char *query);
 extern void kill_work(void);
 extern void switch_pools(struct pool *selected);
+extern void remove_pool(struct pool *pool);
 extern void write_config(FILE *fcfg);
 extern void log_curses(int prio, const char *f, va_list ap);
 extern void clear_logwin(void);
@@ -657,5 +676,6 @@ extern void tq_freeze(struct thread_q *tq);
 extern void tq_thaw(struct thread_q *tq);
 extern bool successful_connect;
 extern void adl(void);
+extern void app_restart(void);
 
 #endif /* __MINER_H__ */

miner.php

@@ -1,9 +1,11 @@
 <?php
 session_start();
 #
-global $miner, $port, $readonly, $notify;
-$miner = '127.0.0.1'; # hostname or IP address
-$port = 4028;
+global $miner, $port, $readonly, $notify, $rigs;
+#
+# Don't touch these 2 - see $rigs below
+$miner = null;
+$port = null;
 #
 # Set $readonly to true to force miner.php to be readonly
 # Set $readonly to false then it will check cgminer 'privileged'
@@ -15,12 +17,27 @@ $readonly = false;
 #  coz it doesn't have notify - it just shows the error status table
 $notify = true;
 #
+# Set $rigs to an array of your cgminer rigs that are running
+#  format: 'IP:Port' or 'Host:Port'
+# If you only have one rig, it will just show the detail of that rig
+# If you have more than one rig it will show a summary of all the rigs
+#  with buttons to show the details of each rig
+# e.g. $rigs = array('127.0.0.1:4028','myrig.com:4028');
+$rigs = array('127.0.0.1:4028');
+#
 $here = $_SERVER['PHP_SELF'];
 #
-function htmlhead()
+global $tablebegin, $tableend, $warnfont, $warnoff;
+$tablebegin = '<tr><td><table border=1 cellpadding=5 cellspacing=0>';
+$tableend = '</table></td></tr>';
+$warnfont = '<font color=red><b>';
+$warnoff = '</b></font>';
+
+#
+function htmlhead($checkapi)
 {
  global $error, $readonly, $here;
- if ($readonly === false)
+ if ($readonly === false && $checkapi === true)
  {
 	$access = api('privileged');
 	if ($error != null
@@ -36,17 +53,18 @@ td.h { color:blue; font-family:verdana,arial,sans; font-size:13pt; background:#d
 td.err { color:black; font-family:verdana,arial,sans; font-size:13pt; background:#ff3050 }
 td.warn { color:black; font-family:verdana,arial,sans; font-size:13pt; background:#ffb050 }
 td.sta { color:green; font-family:verdana,arial,sans; font-size:13pt; }
+td.tot { color:blue; font-family:verdana,arial,sans; font-size:13pt; background:#fff8f2 }
 </style>
 </head><body bgcolor=#ecffff>
 <script type='text/javascript'>
 function pr(a,m){if(m!=null){if(!confirm(m+'?'))return}window.location="<?php echo $here ?>"+a}
 <?php
- if ($readonly === false)
+ if ($readonly === false && $checkapi === true)
  {
 ?>
 function prc(a,m){pr('?arg='+a,m)}
-function prs(a){var c=a.substr(3);var z=c.split('|',2);var m=z[0].substr(0,1).toUpperCase()+z[0].substr(1)+' GPU '+z[1];prc(a,m)}
-function prs2(a,n){var v=document.getElementById('gi'+n).value;var c=a.substr(3);var z=c.split('|',2);var m='Set GPU '+z[1]+' '+z[0].substr(0,1).toUpperCase()+z[0].substr(1)+' to '+v;prc(a+','+v,m)}
+function prs(a,r){var c=a.substr(3);var z=c.split('|',2);var m=z[0].substr(0,1).toUpperCase()+z[0].substr(1)+' GPU '+z[1];prc(a+'&rig='+r,m)}
+function prs2(a,n,r){var v=document.getElementById('gi'+n).value;var c=a.substr(3);var z=c.split('|',2);var m='Set GPU '+z[1]+' '+z[0].substr(0,1).toUpperCase()+z[0].substr(1)+' to '+v;prc(a+','+v+'&rig='+r,m)}
 <?php
  }
 ?>
@@ -219,7 +237,14 @@ function fmt($section, $name, $value)
 			if ($value == 0)
 				$ret = sprintf("%dh$b%02dm$b%02ds", $h, $m, $s);
 			else
-				$ret = sprintf("%ddays$b%02dh$b%02dm$b%02ds", $value, $h, $m, $s);
+			{
+				if ($value == 1)
+					$days = '';
+				else
+					$days = 's';
+
+				$ret = sprintf("%dday$days$b%02dh$b%02dm$b%02ds", $value, $h, $m, $s);
+			}
 		}
 	}
 	break;
@@ -250,10 +275,56 @@ function fmt($section, $name, $value)
  case 'SUMMARY.Utility':
 	$ret = $value.'/m';
 	break;
- case 'GPU.Temperature':
  case 'PGA.Temperature':
 	$ret = $value.'&deg;C';
 	break;
+ case 'GPU.Temperature':
+	$ret = $value.'&deg;C';
+ case 'GPU.Fan Speed':
+ case 'GPU.Fan Percent':
+ case 'GPU.GPU Clock':
+ case 'GPU.Memory Clock':
+ case 'GPU.GPU Voltage':
+ case 'GPU.GPU Activity':
+	if ($value == 0)
+		$class = $warnclass;
+	break;
+ case 'GPU.MHS av':
+ case 'PGA.MHS av':
+ case 'SUMMARY.MHS av':
+ case 'GPU.Total MH':
+ case 'PGA.Total MH':
+ case 'SUMMARY.Total MH':
+ case 'SUMMARY.Getworks':
+ case 'GPU.Accepted':
+ case 'PGA.Accepted':
+ case 'SUMMARY.Accepted':
+ case 'GPU.Rejected':
+ case 'PGA.Rejected':
+ case 'SUMMARY.Rejected':
+ case 'SUMMARY.Local Work':
+ case 'POOL.Getworks':
+ case 'POOL.Accepted':
+ case 'POOL.Rejected':
+ case 'POOL.Discarded':
+	$parts = explode('.', $value, 2);
+	if (count($parts) == 1)
+		$dec = '';
+	else
+		$dec = '.'.$parts[1];
+	$ret = number_format($parts[0]).$dec;
+	break;
+ case 'GPU.Status':
+ case 'PGA.Status':
+ case 'POOL.Status':
+	if ($value != 'Alive')
+		$class = $errorclass;
+	break;
+ case 'GPU.Enabled':
+ case 'PGA.Enabled':
+	if ($value != 'Y')
+		$class = $warnclass;
+	break;
  }
 
  if ($section == 'NOTIFY' && substr($name, 0, 1) == '*' && $value != '0')
@@ -287,22 +358,20 @@ function showhead($cmd, $item, $values)
  echo '</tr>';
 }
 #
-function details($cmd, $list)
+function details($cmd, $list, $rig)
 {
+ global $tablebegin, $tableend;
  global $poolcmd, $readonly;
 
  $dfmt = 'H:i:s j-M-Y \U\T\CP';
 
  $stas = array('S' => 'Success', 'W' => 'Warning', 'I' => 'Informational', 'E' => 'Error', 'F' => 'Fatal');
 
- $tb = '<tr><td><table border=1 cellpadding=5 cellspacing=0>';
- $te = '</table></td></tr>';
-
- echo $tb;
+ echo $tablebegin;
 
  echo '<tr><td class=sta>Date: '.date($dfmt).'</td></tr>';
 
- echo $te.$tb;
+ echo $tableend.$tablebegin;
 
  if (isset($list['STATUS']))
  {
@@ -328,7 +397,7 @@ function details($cmd, $list)
 
 	if ($sectionname != $section)
 	{
-		echo $te.$tb;
+		echo $tableend.$tablebegin;
 		showhead($cmd, $item, $values);
 		$section = $sectionname;
 	}
@@ -353,7 +422,7 @@ function details($cmd, $list)
 			else
 			{
 				echo "<input type=button value='Pool $pool'";
-				echo " onclick='prc(\"$pcmd|$pool\",\"$name Pool $pool\")'>";
+				echo " onclick='prc(\"$pcmd|$pool&rig=$rig\",\"$name Pool $pool\")'>";
 			}
 			echo '</td>';
 		}
@@ -361,14 +430,16 @@ function details($cmd, $list)
 
 	echo '</tr>';
  }
- echo $te;
+
+ echo $tableend;
 }
 #
 global $devs;
 $devs = null;
 #
-function gpubuttons($count)
+function gpubuttons($count, $rig)
 {
+ global $tablebegin, $tableend;
  global $devs;
 
  $basic = array( 'GPU', 'Enable', 'Disable', 'Restart' );
@@ -379,10 +450,7 @@ function gpubuttons($count)
 			'mem' => 'Memory Clock',
 			'vddc' => 'GPU Voltage' );
 
- $tb = '<tr><td><table border=1 cellpadding=5 cellspacing=0>';
- $te = '</table></td></tr>';
-
- echo $tb.'<tr>';
+ echo $tablebegin.'<tr>';
 
  foreach ($basic as $head)
 	echo "<td>$head</td>";
@@ -405,7 +473,7 @@ function gpubuttons($count)
 		{
 			echo "<input type=button value='$name $c' onclick='prs(\"gpu";
 			echo strtolower($name);
-			echo "|$c\")'>";
+			echo "|$c\",$rig)'>";
 		}
 
 		echo '</td>';
@@ -419,7 +487,7 @@ function gpubuttons($count)
 		else
 		{
 			$value = $devs["GPU$c"][$des];
-			echo "<input type=button value='Set $c:' onclick='prs2(\"gpu$name|$c\",$n)'>";
+			echo "<input type=button value='Set $c:' onclick='prs2(\"gpu$name|$c\",$n,$rig)'>";
 			echo "<input size=7 type=text name=gi$n value='$value' id=gi$n>";
 			$n++;
 		}
@@ -429,35 +497,37 @@ function gpubuttons($count)
 
  }
 
- echo '</tr>'.$te;
+ echo '</tr>'.$tableend;
 }
 #
-function processgpus($rd, $ro)
+function processgpus($rig)
 {
  global $error;
+ global $warnfont, $warnoff;
 
  $gpus = api('gpucount');
 
  if ($error != null)
-	echo '<tr><td>Error getting GPU count: '.$rd.$error.$ro.'</td></tr>';
+	echo '<tr><td>Error getting GPU count: '.$warnfont.$error.$warnoff.'</td></tr>';
  else
  {
 	if (!isset($gpus['GPUS']['Count']))
-		echo '<tr><td>No GPU count returned: '.$rd.$gpus['STATUS']['STATUS'].' '.$gpus['STATUS']['Msg'].$ro.'</td></tr>';
+		echo '<tr><td>No GPU count returned: '.$warnfont.$gpus['STATUS']['STATUS'].' '.$gpus['STATUS']['Msg'].$ro.'</td></tr>';
 	else
 	{
 		$count = $gpus['GPUS']['Count'];
 		if ($count == 0)
 			echo '<tr><td>No GPUs</td></tr>';
 		else
-			gpubuttons($count);
+			gpubuttons($count, $rig);
 	}
  }
 }
 #
-function process($cmds, $rd, $ro)
+function process($cmds, $rig)
 {
  global $error, $devs;
+ global $warnfont, $warnoff;
 
  foreach ($cmds as $cmd => $des)
  {
@@ -465,13 +535,13 @@ function process($cmds, $rd, $ro)
 
 	if ($error != null)
 	{
-		echo "<tr><td>Error getting $des: ";
-		echo $rd.$error.$ro.'</td></tr>';
+		echo "<tr><td colspan=100>Error getting $des: ";
+		echo $warnfont.$error.$warnoff.'</td></tr>';
 		break;
 	}
 	else
 	{
-		details($cmd, $process);
+		details($cmd, $process, $rig);
 		echo '<tr><td><br><br></td></tr>';
 		if ($cmd == 'devs')
 			$devs = $process;
@@ -479,25 +549,152 @@ function process($cmds, $rd, $ro)
  }
 }
 #
-function display()
+# $head is a hack but this is just a demo anyway :)
+function doforeach($cmd, $des, $sum, $head)
+{
+ global $miner, $port;
+ global $error, $readonly, $notify, $rigs;
+ global $tablebegin, $tableend, $warnfont, $warnoff;
+
+ $header = $head;
+ $anss = array();
+
+ $count = 0;
+ foreach ($rigs as $rig)
+ {
+	$parts = explode(':', $rig, 2);
+	if (count($parts) == 2)
+	{
+		$miner = $parts[0];
+		$port = $parts[1];
+
+		$ans = api($cmd);
+
+		if ($error != null)
+		{
+			echo "<tr><td colspan=100>Error on rig $count getting $des: ";
+			echo $warnfont.$error.$warnoff.'</td></tr>';
+			$error = null;
+		}
+		else
+			$anss[$count] = $ans;
+	}
+	$count++;
+ }
+
+ if (count($anss) == 0)
+ {
+	echo "<tr><td>Failed to access any rigs successfully</td></tr>";
+	return;
+ }
+
+ $total = array();
+
+ foreach ($anss as $rig => $ans)
+ {
+	foreach ($ans as $item => $row)
+	{
+		if ($item == 'STATUS')
+			continue;
+
+		if (count($row) > count($header))
+		{
+			$header = $head;
+			foreach ($row as $name => $value)
+				if (!isset($header[$name]))
+					$header[$name] = '';
+		}
+
+		if ($sum != null)
+			foreach ($sum as $name)
+			{
+				if (isset($row[$name]))
+				{
+					if (isset($total[$name]))
+						$total[$name] += $row[$name];
+					else
+						$total[$name] = $row[$name];
+				}
+			}
+	}
+ }
+
+ if ($sum != null)
+	$anss['total']['total'] = $total;
+
+ showhead('', null, $header);
+
+ $section = '';
+
+ foreach ($anss as $rig => $ans)
+ {
+	foreach ($ans as $item => $row)
+	{
+		if ($item == 'STATUS')
+			continue;
+
+		echo '<tr>';
+
+		$newsection = preg_replace('/\d/', '', $item);
+		if ($newsection != 'total')
+			$section = $newsection;
+
+		foreach ($header as $name => $x)
+		{
+			if ($name == '')
+			{
+				if ($rig === 'total')
+					echo "<td align=right class=tot>Total:</td>";
+				else
+					echo "<td align=right><input type=button value='Rig $rig' onclick='pr(\"?rig=$rig\",null)'></td>";
+			}
+			else
+			{
+				if (isset($row[$name]))
+					list($showvalue, $class) = fmt($section, $name, $row[$name]);
+				else
+				{
+					$class = '';
+					$showvalue = '&nbsp;';
+				}
+
+				if ($rig === 'total' and $class == '')
+					$class = ' class=tot';
+
+				echo "<td$class align=right>$showvalue</td>";
+			}
+		}
+
+		echo '</tr>';
+	}
+ }
+}
+#
+function doOne($rig, $preprocess)
 {
  global $error, $readonly, $notify;
+ global $rigs;
 
- $error = null;
+ htmlhead(true);
 
- $rd = '<font color=red><b>';
- $ro = '</b></font>';
+ $error = null;
 
  echo "<tr><td><table cellpadding=0 cellspacing=0 border=0><tr><td>";
- echo "<input type=button value='Refresh' onclick='pr(\"\",null)'>";
- echo "</td><td width=100%>&nbsp;</td><td>";
+ echo "<input type=button value='Refresh' onclick='pr(\"?rig=$rig\",null)'></td>";
+ if (count($rigs) > 1)
+	echo "<td><input type=button value='Summary' onclick='pr(\"\",null)'></td>";
+ echo "<td width=100%>&nbsp;</td><td>";
  if ($readonly === false)
-	echo "<input type=button value='Quit' onclick='prc(\"quit\",\"Quit CGMiner\")'>";
+ {
+	$msg = 'Quit CGMiner';
+	if (count($rigs) > 1)
+		$msg .= " Rig $rig";
+	echo "<input type=button value='Quit' onclick='prc(\"quit&rig=$rig\",\"$msg\")'>";
+ }
  echo "</td></tr></table></td></tr>";
 
- $arg = trim(getparam('arg', true));
- if ($arg != null and $arg != '')
-	process(array($arg => $arg), $rd, $ro);
+ if ($preprocess != null)
+	process(array($preprocess => $preprocess), $rig);
 
  $cmds = array(	'devs'    => 'device list',
 		'summary' => 'summary information',
@@ -508,13 +705,108 @@ function display()
 
  $cmds['config'] = 'cgminer config';
 
- process($cmds, $rd, $ro);
+ process($cmds, $rig);
 
  if ($error == null && $readonly === false)
-	processgpus($rd, $ro);
+	processgpus($rig);
+}
+#
+function display()
+{
+ global $tablebegin, $tableend;
+ global $miner, $port;
+ global $error, $readonly, $notify, $rigs;
+
+ $rig = trim(getparam('rig', true));
+
+ $arg = trim(getparam('arg', true));
+ $preprocess = null;
+ if ($arg != null and $arg != '')
+ {
+	$num = null;
+	if ($rig != null and $rig != '')
+	{
+		if ($rig >= 0 and $rig < count($rigs))
+			$num = $rig;
+	}
+	else
+		if (count($rigs) == 0)
+			$num = 0;
+
+	if ($num != null)
+	{
+		$parts = explode(':', $rigs[$num], 2);
+		if (count($parts) == 2)
+		{
+			$miner = $parts[0];
+			$port = $parts[1];
+
+			$preprocess = $arg;
+		}
+	}
+ }
+
+ if ($rigs == null or count($rigs) == 0)
+ {
+	echo "<tr><td>No rigs defined</td></tr>";
+	return;
+ }
+
+ if (count($rigs) == 1)
+ {
+	$parts = explode(':', $rigs[0], 2);
+	if (count($parts) == 2)
+	{
+		$miner = $parts[0];
+		$port = $parts[1];
+
+		doOne(0, $preprocess);
+	}
+	else
+		echo '<tr><td>Invalid "$rigs" array</td></tr>';
+
+	return;
+ }
+
+ if ($rig != null and $rig != '' and $rig >= 0 and $rig < count($rigs))
+ {
+	$parts = explode(':', $rigs[$rig], 2);
+	if (count($parts) == 2)
+	{
+		$miner = $parts[0];
+		$port = $parts[1];
+
+		doOne($rig, $preprocess);
+	}
+	else
+		echo '<tr><td>Invalid "$rigs" array</td></tr>';
+
+	return;
+ }
+
+ htmlhead(false);
+
+ echo "<tr><td><table cellpadding=0 cellspacing=0 border=0><tr><td>";
+ echo "<input type=button value='Refresh' onclick='pr(\"\",null)'>";
+ echo "</td></tr></table></td></tr>";
+
+ if ($preprocess != null)
+	process(array($preprocess => $preprocess), $rig);
+
+ echo $tablebegin;
+ $sum = array('MHS av', 'Getworks', 'Found Blocks', 'Accepted', 'Rejected', 'Discarded', 'Stale', 'Utility', 'Local Work', 'Total MH');
+ doforeach('summary', 'summary information', $sum, array());
+ echo $tableend;
+ echo '<tr><td><br><br></td></tr>';
+ echo $tablebegin;
+ doforeach('devs', 'device list', $sum, array(''=>'','ID'=>'','Name'=>''));
+ echo $tableend;
+ echo '<tr><td><br><br></td></tr>';
+ echo $tablebegin;
+ doforeach('pools', 'pool list', $sum, array(''=>''));
+ echo $tableend;
 }
 #
-htmlhead();
 display();
 #
 ?>

mkinstalldirs

@@ -91,1262 +91,1263 @@ void search(const uint state0, const uint state1, const uint state2, const uint
 	const u nonce = base + (uint)(get_global_id(0));
 #endif
 
-Vals[0]=Preval0;
-Vals[0]+=nonce;
+Vals[5]=Preval0;
+Vals[5]+=nonce;
 
-Vals[3]=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
-Vals[3]+=ch(Vals[0],b1,c1);
-Vals[3]+=D1A;
+Vals[0]=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
+Vals[0]+=ch(Vals[5],b1,c1);
+Vals[0]+=D1A;
 
-Vals[7]=Vals[3];
-Vals[7]+=h1;
+Vals[2]=Vals[0];
+Vals[2]+=h1;
 
-Vals[4]=PreVal4addT1;
-Vals[4]+=nonce;
-Vals[3]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
-
-Vals[2]=C1addK5;
-Vals[2]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
-Vals[2]+=ch(Vals[7],Vals[0],b1);
+Vals[1]=PreVal4addT1;
+Vals[1]+=nonce;
+Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
 
-Vals[6]=Vals[2];
-Vals[6]+=g1;
-Vals[3]+=Ma2(g1,Vals[4],f1);
-Vals[2]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
-Vals[2]+=Ma2(f1,Vals[3],Vals[4]);
+Vals[6]=C1addK5;
+Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
+Vals[6]+=ch(Vals[2],Vals[5],b1);
 
-Vals[1]=B1addK6;
-Vals[1]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
-Vals[1]+=ch(Vals[6],Vals[7],Vals[0]);
+Vals[3]=Vals[6];
+Vals[3]+=g1;
+Vals[0]+=Ma2(g1,Vals[1],f1);
+Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
+Vals[6]+=Ma2(f1,Vals[0],Vals[1]);
 
-Vals[5]=Vals[1];
-Vals[5]+=f1;
+Vals[7]=B1addK6;
+Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
+Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
 
-Vals[1]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
-Vals[1]+=Ma(Vals[4],Vals[2],Vals[3]);
+Vals[4]=Vals[7];
+Vals[4]+=f1;
 
-Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
-Vals[0]+=ch(Vals[5],Vals[6],Vals[7]);
-Vals[0]+=K[7];
-Vals[4]+=Vals[0];
-Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
-Vals[0]+=Ma(Vals[3],Vals[1],Vals[2]);
+Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
+Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
 
-Vals[7]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
-Vals[7]+=ch(Vals[4],Vals[5],Vals[6]);
-Vals[7]+=K[8];
-Vals[3]+=Vals[7];
-Vals[7]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
-Vals[7]+=Ma(Vals[2],Vals[0],Vals[1]);
-
-Vals[6]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
-Vals[6]+=ch(Vals[3],Vals[4],Vals[5]);
-Vals[6]+=K[9];
-Vals[2]+=Vals[6];
-Vals[6]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
-Vals[6]+=Ma(Vals[1],Vals[7],Vals[0]);
-
-Vals[5]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
-Vals[5]+=ch(Vals[2],Vals[3],Vals[4]);
-Vals[5]+=K[10];
+Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
+Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
+Vals[5]+=K[7];
 Vals[1]+=Vals[5];
-Vals[5]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
-Vals[5]+=Ma(Vals[0],Vals[6],Vals[7]);
+Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
+Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
 
-Vals[4]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
-Vals[4]+=ch(Vals[1],Vals[2],Vals[3]);
-Vals[4]+=K[11];
-Vals[0]+=Vals[4];
-Vals[4]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
-Vals[4]+=Ma(Vals[7],Vals[5],Vals[6]);
+Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
+Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
+Vals[2]+=K[8];
+Vals[0]+=Vals[2];
+Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
+Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
 
 Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
-Vals[3]+=ch(Vals[0],Vals[1],Vals[2]);
-Vals[3]+=K[12];
-Vals[7]+=Vals[3];
-Vals[3]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
-Vals[3]+=Ma(Vals[6],Vals[4],Vals[5]);
-
-Vals[2]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
-Vals[2]+=ch(Vals[7],Vals[0],Vals[1]);
-Vals[2]+=K[13];
-Vals[6]+=Vals[2];
-Vals[2]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
-Vals[2]+=Ma(Vals[5],Vals[3],Vals[4]);
-
-Vals[1]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
-Vals[1]+=ch(Vals[6],Vals[7],Vals[0]);
-Vals[1]+=K[14];
+Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
+Vals[3]+=K[9];
+Vals[6]+=Vals[3];
+Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
+Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
+
+Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
+Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
+Vals[4]+=K[10];
+Vals[7]+=Vals[4];
+Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
+Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
+
+Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
+Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
+Vals[1]+=K[11];
 Vals[5]+=Vals[1];
-Vals[1]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
-Vals[1]+=Ma(Vals[4],Vals[2],Vals[3]);
+Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
+Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
 
 Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
-Vals[0]+=ch(Vals[5],Vals[6],Vals[7]);
-Vals[0]+=0xC19BF3F4U;
-Vals[4]+=Vals[0];
+Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
+Vals[0]+=K[12];
+Vals[2]+=Vals[0];
 Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
-Vals[0]+=Ma(Vals[3],Vals[1],Vals[2]);
+Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
+
+Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
+Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
+Vals[6]+=K[13];
+Vals[3]+=Vals[6];
+Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
+Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
+
+Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
+Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
+Vals[7]+=K[14];
+Vals[4]+=Vals[7];
+Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
+Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
+
+Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
+Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
+Vals[5]+=0xC19BF3F4U;
+Vals[1]+=Vals[5];
+Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
+Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
 
-Vals[7]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
-Vals[7]+=ch(Vals[4],Vals[5],Vals[6]);
-Vals[7]+=W16addK16;
-Vals[3]+=Vals[7];
-Vals[7]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
-Vals[7]+=Ma(Vals[2],Vals[0],Vals[1]);
+Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
+Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
+Vals[2]+=W16addK16;
+Vals[0]+=Vals[2];
+Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
+Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
 
-Vals[6]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
-Vals[6]+=ch(Vals[3],Vals[4],Vals[5]);
-Vals[6]+=W17addK17;
-Vals[2]+=Vals[6];
-Vals[6]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
-Vals[6]+=Ma(Vals[1],Vals[7],Vals[0]);
+Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
+Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
+Vals[3]+=W17addK17;
+Vals[6]+=Vals[3];
+Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
+Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
 
 W[2]=(rotr(nonce,7)^rotr(nonce,18)^(nonce>>3U));
 W[2]+=fw2;
-Vals[5]+=W[2];
-Vals[5]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
-Vals[5]+=ch(Vals[2],Vals[3],Vals[4]);
-Vals[5]+=K[18];
-Vals[1]+=Vals[5];
-Vals[5]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
-Vals[5]+=Ma(Vals[0],Vals[6],Vals[7]);
+Vals[4]+=W[2];
+Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
+Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
+Vals[4]+=K[18];
+Vals[7]+=Vals[4];
+Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
+Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
 
 W[3]=nonce;
 W[3]+=fw3;
-Vals[4]+=W[3];
-Vals[4]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
-Vals[4]+=ch(Vals[1],Vals[2],Vals[3]);
-Vals[4]+=K[19];
-Vals[0]+=Vals[4];
-Vals[4]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
-Vals[4]+=Ma(Vals[7],Vals[5],Vals[6]);
+Vals[1]+=W[3];
+Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
+Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
+Vals[1]+=K[19];
+Vals[5]+=Vals[1];
+Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
+Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
 
 W[4]=(rotr(W[2],17)^rotr(W[2],19)^(W[2]>>10U));
 W[4]+=0x80000000U;
-Vals[3]+=W[4];
-Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
-Vals[3]+=ch(Vals[0],Vals[1],Vals[2]);
-Vals[3]+=K[20];
-Vals[7]+=Vals[3];
-Vals[3]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
-Vals[3]+=Ma(Vals[6],Vals[4],Vals[5]);
+Vals[0]+=W[4];
+Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
+Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
+Vals[0]+=K[20];
+Vals[2]+=Vals[0];
+Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
+Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
 
 W[5]=(rotr(W[3],17)^rotr(W[3],19)^(W[3]>>10U));
-Vals[2]+=W[5];
-Vals[2]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
-Vals[2]+=ch(Vals[7],Vals[0],Vals[1]);
-Vals[2]+=K[21];
-Vals[6]+=Vals[2];
-Vals[2]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
-Vals[2]+=Ma(Vals[5],Vals[3],Vals[4]);
+Vals[6]+=W[5];
+Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
+Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
+Vals[6]+=K[21];
+Vals[3]+=Vals[6];
+Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
+Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
 
 W[6]=(rotr(W[4],17)^rotr(W[4],19)^(W[4]>>10U));
 W[6]+=0x00000280U;
-Vals[1]+=W[6];
-Vals[1]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
-Vals[1]+=ch(Vals[6],Vals[7],Vals[0]);
-Vals[1]+=K[22];
-Vals[5]+=Vals[1];
-Vals[1]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
-Vals[1]+=Ma(Vals[4],Vals[2],Vals[3]);
+Vals[7]+=W[6];
+Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
+Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
+Vals[7]+=K[22];
+Vals[4]+=Vals[7];
+Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
+Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
 
 W[7]=(rotr(W[5],17)^rotr(W[5],19)^(W[5]>>10U));
 W[7]+=fw0;
-Vals[0]+=W[7];
-Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
-Vals[0]+=ch(Vals[5],Vals[6],Vals[7]);
-Vals[0]+=K[23];
-Vals[4]+=Vals[0];
-Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
-Vals[0]+=Ma(Vals[3],Vals[1],Vals[2]);
+Vals[5]+=W[7];
+Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
+Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
+Vals[5]+=K[23];
+Vals[1]+=Vals[5];
+Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
+Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
 
 W[8]=(rotr(W[6],17)^rotr(W[6],19)^(W[6]>>10U));
 W[8]+=fw1;
-Vals[7]+=W[8];
-Vals[7]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
-Vals[7]+=ch(Vals[4],Vals[5],Vals[6]);
-Vals[7]+=K[24];
-Vals[3]+=Vals[7];
-Vals[7]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
-Vals[7]+=Ma(Vals[2],Vals[0],Vals[1]);
+Vals[2]+=W[8];
+Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
+Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
+Vals[2]+=K[24];
+Vals[0]+=Vals[2];
+Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
+Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
 
 W[9]=W[2];
 W[9]+=(rotr(W[7],17)^rotr(W[7],19)^(W[7]>>10U));
-Vals[6]+=W[9];
-Vals[6]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
-Vals[6]+=ch(Vals[3],Vals[4],Vals[5]);
-Vals[6]+=K[25];
-Vals[2]+=Vals[6];
-Vals[6]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
-Vals[6]+=Ma(Vals[1],Vals[7],Vals[0]);
+Vals[3]+=W[9];
+Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
+Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
+Vals[3]+=K[25];
+Vals[6]+=Vals[3];
+Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
+Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
 
 W[10]=W[3];
 W[10]+=(rotr(W[8],17)^rotr(W[8],19)^(W[8]>>10U));
-Vals[5]+=W[10];
-Vals[5]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
-Vals[5]+=ch(Vals[2],Vals[3],Vals[4]);
-Vals[5]+=K[26];
-Vals[1]+=Vals[5];
-Vals[5]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
-Vals[5]+=Ma(Vals[0],Vals[6],Vals[7]);
+Vals[4]+=W[10];
+Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
+Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
+Vals[4]+=K[26];
+Vals[7]+=Vals[4];
+Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
+Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
 
 W[11]=W[4];
 W[11]+=(rotr(W[9],17)^rotr(W[9],19)^(W[9]>>10U));
-Vals[4]+=W[11];
-Vals[4]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
-Vals[4]+=ch(Vals[1],Vals[2],Vals[3]);
-Vals[4]+=K[27];
-Vals[0]+=Vals[4];
-Vals[4]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
-Vals[4]+=Ma(Vals[7],Vals[5],Vals[6]);
+Vals[1]+=W[11];
+Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
+Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
+Vals[1]+=K[27];
+Vals[5]+=Vals[1];
+Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
+Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
 
 W[12]=W[5];
 W[12]+=(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U));
-Vals[3]+=W[12];
-Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
-Vals[3]+=ch(Vals[0],Vals[1],Vals[2]);
-Vals[3]+=K[28];
-Vals[7]+=Vals[3];
-Vals[3]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
-Vals[3]+=Ma(Vals[6],Vals[4],Vals[5]);
+Vals[0]+=W[12];
+Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
+Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
+Vals[0]+=K[28];
+Vals[2]+=Vals[0];
+Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
+Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
 
 W[13]=W[6];
 W[13]+=(rotr(W[11],17)^rotr(W[11],19)^(W[11]>>10U));
-Vals[2]+=W[13];
-Vals[2]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
-Vals[2]+=ch(Vals[7],Vals[0],Vals[1]);
-Vals[2]+=K[29];
-Vals[6]+=Vals[2];
-Vals[2]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
-Vals[2]+=Ma(Vals[5],Vals[3],Vals[4]);
+Vals[6]+=W[13];
+Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
+Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
+Vals[6]+=K[29];
+Vals[3]+=Vals[6];
+Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
+Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
 
 W[14]=0x00a00055U;
 W[14]+=W[7];
 W[14]+=(rotr(W[12],17)^rotr(W[12],19)^(W[12]>>10U));
-Vals[1]+=W[14];
-Vals[1]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
-Vals[1]+=ch(Vals[6],Vals[7],Vals[0]);
-Vals[1]+=K[30];
-Vals[5]+=Vals[1];
-Vals[1]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
-Vals[1]+=Ma(Vals[4],Vals[2],Vals[3]);
+Vals[7]+=W[14];
+Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
+Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
+Vals[7]+=K[30];
+Vals[4]+=Vals[7];
+Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
+Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
 
 W[15]=fw15;
 W[15]+=W[8];
 W[15]+=(rotr(W[13],17)^rotr(W[13],19)^(W[13]>>10U));
-Vals[0]+=W[15];
-Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
-Vals[0]+=ch(Vals[5],Vals[6],Vals[7]);
-Vals[0]+=K[31];
-Vals[4]+=Vals[0];
-Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
-Vals[0]+=Ma(Vals[3],Vals[1],Vals[2]);
+Vals[5]+=W[15];
+Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
+Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
+Vals[5]+=K[31];
+Vals[1]+=Vals[5];
+Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
+Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
 
 W[0]=fw01r;
 W[0]+=W[9];
 W[0]+=(rotr(W[14],17)^rotr(W[14],19)^(W[14]>>10U));
-Vals[7]+=W[0];
-Vals[7]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
-Vals[7]+=ch(Vals[4],Vals[5],Vals[6]);
-Vals[7]+=K[32];
-Vals[3]+=Vals[7];
-Vals[7]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
-Vals[7]+=Ma(Vals[2],Vals[0],Vals[1]);
+Vals[2]+=W[0];
+Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
+Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
+Vals[2]+=K[32];
+Vals[0]+=Vals[2];
+Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
+Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
 
 W[1]=fw1;
 W[1]+=(rotr(W[2],7)^rotr(W[2],18)^(W[2]>>3U));
 W[1]+=W[10];
 W[1]+=(rotr(W[15],17)^rotr(W[15],19)^(W[15]>>10U));
-Vals[6]+=W[1];
-Vals[6]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
-Vals[6]+=ch(Vals[3],Vals[4],Vals[5]);
-Vals[6]+=K[33];
-Vals[2]+=Vals[6];
-Vals[6]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
-Vals[6]+=Ma(Vals[1],Vals[7],Vals[0]);
+Vals[3]+=W[1];
+Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
+Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
+Vals[3]+=K[33];
+Vals[6]+=Vals[3];
+Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
+Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
 
 W[2]+=(rotr(W[3],7)^rotr(W[3],18)^(W[3]>>3U));
 W[2]+=W[11];
 W[2]+=(rotr(W[0],17)^rotr(W[0],19)^(W[0]>>10U));
-Vals[5]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
-Vals[5]+=ch(Vals[2],Vals[3],Vals[4]);
-Vals[5]+=K[34];
-Vals[5]+=W[2];
-Vals[1]+=Vals[5];
-Vals[5]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
-Vals[5]+=Ma(Vals[0],Vals[6],Vals[7]);
+Vals[4]+=W[2];
+Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
+Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
+Vals[4]+=K[34];
+Vals[7]+=Vals[4];
+Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
+Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
 
 W[3]+=(rotr(W[4],7)^rotr(W[4],18)^(W[4]>>3U));
 W[3]+=W[12];
 W[3]+=(rotr(W[1],17)^rotr(W[1],19)^(W[1]>>10U));
-Vals[4]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
-Vals[4]+=ch(Vals[1],Vals[2],Vals[3]);
-Vals[4]+=K[35];
-Vals[4]+=W[3];
-Vals[0]+=Vals[4];
-Vals[4]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
-Vals[4]+=Ma(Vals[7],Vals[5],Vals[6]);
+Vals[1]+=W[3];
+Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
+Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
+Vals[1]+=K[35];
+Vals[5]+=Vals[1];
+Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
+Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
 
 W[4]+=(rotr(W[5],7)^rotr(W[5],18)^(W[5]>>3U));
 W[4]+=W[13];
 W[4]+=(rotr(W[2],17)^rotr(W[2],19)^(W[2]>>10U));
-Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
-Vals[3]+=ch(Vals[0],Vals[1],Vals[2]);
-Vals[3]+=K[36];
-Vals[3]+=W[4];
-Vals[7]+=Vals[3];
-Vals[3]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
-Vals[3]+=Ma(Vals[6],Vals[4],Vals[5]);
+Vals[0]+=W[4];
+Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
+Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
+Vals[0]+=K[36];
+Vals[2]+=Vals[0];
+Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
+Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
 
 W[5]+=(rotr(W[6],7)^rotr(W[6],18)^(W[6]>>3U));
 W[5]+=W[14];
 W[5]+=(rotr(W[3],17)^rotr(W[3],19)^(W[3]>>10U));
-Vals[2]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
-Vals[2]+=ch(Vals[7],Vals[0],Vals[1]);
-Vals[2]+=K[37];
-Vals[2]+=W[5];
-Vals[6]+=Vals[2];
-Vals[2]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
-Vals[2]+=Ma(Vals[5],Vals[3],Vals[4]);
+Vals[6]+=W[5];
+Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
+Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
+Vals[6]+=K[37];
+Vals[3]+=Vals[6];
+Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
+Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
 
 W[6]+=(rotr(W[7],7)^rotr(W[7],18)^(W[7]>>3U));
 W[6]+=W[15];
 W[6]+=(rotr(W[4],17)^rotr(W[4],19)^(W[4]>>10U));
-Vals[1]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
-Vals[1]+=ch(Vals[6],Vals[7],Vals[0]);
-Vals[1]+=K[38];
-Vals[1]+=W[6];
-Vals[5]+=Vals[1];
-Vals[1]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
-Vals[1]+=Ma(Vals[4],Vals[2],Vals[3]);
+Vals[7]+=W[6];
+Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
+Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
+Vals[7]+=K[38];
+Vals[4]+=Vals[7];
+Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
+Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
 
 W[7]+=(rotr(W[8],7)^rotr(W[8],18)^(W[8]>>3U));
 W[7]+=W[0];
 W[7]+=(rotr(W[5],17)^rotr(W[5],19)^(W[5]>>10U));
-Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
-Vals[0]+=ch(Vals[5],Vals[6],Vals[7]);
-Vals[0]+=K[39];
-Vals[0]+=W[7];
-Vals[4]+=Vals[0];
-Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
-Vals[0]+=Ma(Vals[3],Vals[1],Vals[2]);
+Vals[5]+=W[7];
+Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
+Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
+Vals[5]+=K[39];
+Vals[1]+=Vals[5];
+Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
+Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
 
 W[8]+=(rotr(W[9],7)^rotr(W[9],18)^(W[9]>>3U));
 W[8]+=W[1];
 W[8]+=(rotr(W[6],17)^rotr(W[6],19)^(W[6]>>10U));
-Vals[7]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
-Vals[7]+=ch(Vals[4],Vals[5],Vals[6]);
-Vals[7]+=K[40];
-Vals[7]+=W[8];
-Vals[3]+=Vals[7];
-Vals[7]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
-Vals[7]+=Ma(Vals[2],Vals[0],Vals[1]);
+Vals[2]+=W[8];
+Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
+Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
+Vals[2]+=K[40];
+Vals[0]+=Vals[2];
+Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
+Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
 
 W[9]+=(rotr(W[10],7)^rotr(W[10],18)^(W[10]>>3U));
 W[9]+=W[2];
 W[9]+=(rotr(W[7],17)^rotr(W[7],19)^(W[7]>>10U));
-Vals[6]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
-Vals[6]+=ch(Vals[3],Vals[4],Vals[5]);
-Vals[6]+=K[41];
-Vals[6]+=W[9];
-Vals[2]+=Vals[6];
-Vals[6]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
-Vals[6]+=Ma(Vals[1],Vals[7],Vals[0]);
+Vals[3]+=W[9];
+Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
+Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
+Vals[3]+=K[41];
+Vals[6]+=Vals[3];
+Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
+Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
 
 W[10]+=(rotr(W[11],7)^rotr(W[11],18)^(W[11]>>3U));
 W[10]+=W[3];
 W[10]+=(rotr(W[8],17)^rotr(W[8],19)^(W[8]>>10U));
-Vals[5]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
-Vals[5]+=ch(Vals[2],Vals[3],Vals[4]);
-Vals[5]+=K[42];
-Vals[5]+=W[10];
-Vals[1]+=Vals[5];
-Vals[5]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
-Vals[5]+=Ma(Vals[0],Vals[6],Vals[7]);
+Vals[4]+=W[10];
+Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
+Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
+Vals[4]+=K[42];
+Vals[7]+=Vals[4];
+Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
+Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
 
 W[11]+=(rotr(W[12],7)^rotr(W[12],18)^(W[12]>>3U));
 W[11]+=W[4];
 W[11]+=(rotr(W[9],17)^rotr(W[9],19)^(W[9]>>10U));
-Vals[4]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
-Vals[4]+=ch(Vals[1],Vals[2],Vals[3]);
-Vals[4]+=K[43];
-Vals[4]+=W[11];
-Vals[0]+=Vals[4];
-Vals[4]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
-Vals[4]+=Ma(Vals[7],Vals[5],Vals[6]);
+Vals[1]+=W[11];
+Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
+Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
+Vals[1]+=K[43];
+Vals[5]+=Vals[1];
+Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
+Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
 
 W[12]+=(rotr(W[13],7)^rotr(W[13],18)^(W[13]>>3U));
 W[12]+=W[5];
 W[12]+=(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U));
-Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
-Vals[3]+=ch(Vals[0],Vals[1],Vals[2]);
-Vals[3]+=K[44];
-Vals[3]+=W[12];
-Vals[7]+=Vals[3];
-Vals[3]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
-Vals[3]+=Ma(Vals[6],Vals[4],Vals[5]);
+Vals[0]+=W[12];
+Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
+Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
+Vals[0]+=K[44];
+Vals[2]+=Vals[0];
+Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
+Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
 
 W[13]+=(rotr(W[14],7)^rotr(W[14],18)^(W[14]>>3U));
 W[13]+=W[6];
 W[13]+=(rotr(W[11],17)^rotr(W[11],19)^(W[11]>>10U));
-Vals[2]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
-Vals[2]+=ch(Vals[7],Vals[0],Vals[1]);
-Vals[2]+=K[45];
-Vals[2]+=W[13];
-Vals[6]+=Vals[2];
-Vals[2]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
-Vals[2]+=Ma(Vals[5],Vals[3],Vals[4]);
+Vals[6]+=W[13];
+Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
+Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
+Vals[6]+=K[45];
+Vals[3]+=Vals[6];
+Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
+Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
 
 W[14]+=(rotr(W[15],7)^rotr(W[15],18)^(W[15]>>3U));
 W[14]+=W[7];
 W[14]+=(rotr(W[12],17)^rotr(W[12],19)^(W[12]>>10U));
-Vals[1]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
-Vals[1]+=ch(Vals[6],Vals[7],Vals[0]);
-Vals[1]+=K[46];
-Vals[1]+=W[14];
-Vals[5]+=Vals[1];
-Vals[1]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
-Vals[1]+=Ma(Vals[4],Vals[2],Vals[3]);
+Vals[7]+=W[14];
+Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
+Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
+Vals[7]+=K[46];
+Vals[4]+=Vals[7];
+Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
+Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
 
 W[15]+=(rotr(W[0],7)^rotr(W[0],18)^(W[0]>>3U));
 W[15]+=W[8];
 W[15]+=(rotr(W[13],17)^rotr(W[13],19)^(W[13]>>10U));
-Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
-Vals[0]+=ch(Vals[5],Vals[6],Vals[7]);
-Vals[0]+=K[47];
-Vals[0]+=W[15];
-Vals[4]+=Vals[0];
-Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
-Vals[0]+=Ma(Vals[3],Vals[1],Vals[2]);
+Vals[5]+=W[15];
+Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
+Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
+Vals[5]+=K[47];
+Vals[1]+=Vals[5];
+Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
+Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
 
 W[0]+=(rotr(W[1],7)^rotr(W[1],18)^(W[1]>>3U));
 W[0]+=W[9];
 W[0]+=(rotr(W[14],17)^rotr(W[14],19)^(W[14]>>10U));
-Vals[7]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
-Vals[7]+=ch(Vals[4],Vals[5],Vals[6]);
-Vals[7]+=K[48];
-Vals[7]+=W[0];
-Vals[3]+=Vals[7];
-Vals[7]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
-Vals[7]+=Ma(Vals[2],Vals[0],Vals[1]);
+Vals[2]+=W[0];
+Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
+Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
+Vals[2]+=K[48];
+Vals[0]+=Vals[2];
+Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
+Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
 
 W[1]+=(rotr(W[2],7)^rotr(W[2],18)^(W[2]>>3U));
 W[1]+=W[10];
 W[1]+=(rotr(W[15],17)^rotr(W[15],19)^(W[15]>>10U));
-Vals[6]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
-Vals[6]+=ch(Vals[3],Vals[4],Vals[5]);
-Vals[6]+=K[49];
-Vals[6]+=W[1];
-Vals[2]+=Vals[6];
-Vals[6]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
-Vals[6]+=Ma(Vals[1],Vals[7],Vals[0]);
+Vals[3]+=W[1];
+Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
+Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
+Vals[3]+=K[49];
+Vals[6]+=Vals[3];
+Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
+Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
 
 W[2]+=(rotr(W[3],7)^rotr(W[3],18)^(W[3]>>3U));
 W[2]+=W[11];
 W[2]+=(rotr(W[0],17)^rotr(W[0],19)^(W[0]>>10U));
-Vals[5]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
-Vals[5]+=ch(Vals[2],Vals[3],Vals[4]);
-Vals[5]+=K[50];
-Vals[5]+=W[2];
-Vals[1]+=Vals[5];
-Vals[5]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
-Vals[5]+=Ma(Vals[0],Vals[6],Vals[7]);
+Vals[4]+=W[2];
+Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
+Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
+Vals[4]+=K[50];
+Vals[7]+=Vals[4];
+Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
+Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
 
 W[3]+=(rotr(W[4],7)^rotr(W[4],18)^(W[4]>>3U));
 W[3]+=W[12];
 W[3]+=(rotr(W[1],17)^rotr(W[1],19)^(W[1]>>10U));
-Vals[4]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
-Vals[4]+=ch(Vals[1],Vals[2],Vals[3]);
-Vals[4]+=K[51];
-Vals[4]+=W[3];
-Vals[0]+=Vals[4];
-Vals[4]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
-Vals[4]+=Ma(Vals[7],Vals[5],Vals[6]);
+Vals[1]+=W[3];
+Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
+Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
+Vals[1]+=K[51];
+Vals[5]+=Vals[1];
+Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
+Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
 
 W[4]+=(rotr(W[5],7)^rotr(W[5],18)^(W[5]>>3U));
 W[4]+=W[13];
 W[4]+=(rotr(W[2],17)^rotr(W[2],19)^(W[2]>>10U));
-Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
-Vals[3]+=ch(Vals[0],Vals[1],Vals[2]);
-Vals[3]+=K[52];
-Vals[3]+=W[4];
-Vals[7]+=Vals[3];
-Vals[3]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
-Vals[3]+=Ma(Vals[6],Vals[4],Vals[5]);
+Vals[0]+=W[4];
+Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
+Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
+Vals[0]+=K[52];
+Vals[2]+=Vals[0];
+Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
+Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
 
 W[5]+=(rotr(W[6],7)^rotr(W[6],18)^(W[6]>>3U));
 W[5]+=W[14];
 W[5]+=(rotr(W[3],17)^rotr(W[3],19)^(W[3]>>10U));
-Vals[2]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
-Vals[2]+=ch(Vals[7],Vals[0],Vals[1]);
-Vals[2]+=K[53];
-Vals[2]+=W[5];
-Vals[6]+=Vals[2];
-Vals[2]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
-Vals[2]+=Ma(Vals[5],Vals[3],Vals[4]);
+Vals[6]+=W[5];
+Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
+Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
+Vals[6]+=K[53];
+Vals[3]+=Vals[6];
+Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
+Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
 
 W[6]+=(rotr(W[7],7)^rotr(W[7],18)^(W[7]>>3U));
 W[6]+=W[15];
 W[6]+=(rotr(W[4],17)^rotr(W[4],19)^(W[4]>>10U));
-Vals[1]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
-Vals[1]+=ch(Vals[6],Vals[7],Vals[0]);
-Vals[1]+=K[54];
-Vals[1]+=W[6];
-Vals[5]+=Vals[1];
-Vals[1]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
-Vals[1]+=Ma(Vals[4],Vals[2],Vals[3]);
+Vals[7]+=W[6];
+Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
+Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
+Vals[7]+=K[54];
+Vals[4]+=Vals[7];
+Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
+Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
 
 W[7]+=(rotr(W[8],7)^rotr(W[8],18)^(W[8]>>3U));
 W[7]+=W[0];
 W[7]+=(rotr(W[5],17)^rotr(W[5],19)^(W[5]>>10U));
-Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
-Vals[0]+=ch(Vals[5],Vals[6],Vals[7]);
-Vals[0]+=K[55];
-Vals[0]+=W[7];
-Vals[4]+=Vals[0];
-Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
-Vals[0]+=Ma(Vals[3],Vals[1],Vals[2]);
+Vals[5]+=W[7];
+Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
+Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
+Vals[5]+=K[55];
+Vals[1]+=Vals[5];
+Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
+Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
 
 W[8]+=(rotr(W[9],7)^rotr(W[9],18)^(W[9]>>3U));
 W[8]+=W[1];
 W[8]+=(rotr(W[6],17)^rotr(W[6],19)^(W[6]>>10U));
-Vals[7]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
-Vals[7]+=ch(Vals[4],Vals[5],Vals[6]);
-Vals[7]+=K[56];
-Vals[7]+=W[8];
-Vals[3]+=Vals[7];
-Vals[7]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
-Vals[7]+=Ma(Vals[2],Vals[0],Vals[1]);
+Vals[2]+=W[8];
+Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
+Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
+Vals[2]+=K[56];
+Vals[0]+=Vals[2];
+Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
+Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
 
 W[9]+=(rotr(W[10],7)^rotr(W[10],18)^(W[10]>>3U));
 W[9]+=W[2];
 W[9]+=(rotr(W[7],17)^rotr(W[7],19)^(W[7]>>10U));
-Vals[6]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
-Vals[6]+=ch(Vals[3],Vals[4],Vals[5]);
-Vals[6]+=K[57];
-Vals[6]+=W[9];
-Vals[2]+=Vals[6];
-Vals[6]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
-Vals[6]+=Ma(Vals[1],Vals[7],Vals[0]);
+Vals[3]+=W[9];
+Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
+Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
+Vals[3]+=K[57];
+Vals[6]+=Vals[3];
+Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
+Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
 
 W[10]+=(rotr(W[11],7)^rotr(W[11],18)^(W[11]>>3U));
 W[10]+=W[3];
 W[10]+=(rotr(W[8],17)^rotr(W[8],19)^(W[8]>>10U));
-Vals[5]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
-Vals[5]+=ch(Vals[2],Vals[3],Vals[4]);
-Vals[5]+=K[58];
-Vals[5]+=W[10];
-Vals[1]+=Vals[5];
-Vals[5]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
-Vals[5]+=Ma(Vals[0],Vals[6],Vals[7]);
+Vals[4]+=W[10];
+Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
+Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
+Vals[4]+=K[58];
+Vals[7]+=Vals[4];
+Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
+Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
 
 W[11]+=(rotr(W[12],7)^rotr(W[12],18)^(W[12]>>3U));
 W[11]+=W[4];
 W[11]+=(rotr(W[9],17)^rotr(W[9],19)^(W[9]>>10U));
-Vals[4]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
-Vals[4]+=ch(Vals[1],Vals[2],Vals[3]);
-Vals[4]+=K[59];
-Vals[4]+=W[11];
-Vals[0]+=Vals[4];
-Vals[4]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
-Vals[4]+=Ma(Vals[7],Vals[5],Vals[6]);
+Vals[1]+=W[11];
+Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
+Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
+Vals[1]+=K[59];
+Vals[5]+=Vals[1];
+Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
+Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
 
 W[12]+=(rotr(W[13],7)^rotr(W[13],18)^(W[13]>>3U));
 W[12]+=W[5];
 W[12]+=(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U));
-Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
-Vals[3]+=ch(Vals[0],Vals[1],Vals[2]);
-Vals[3]+=K[60];
-Vals[3]+=W[12];
-Vals[7]+=Vals[3];
-Vals[3]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
-Vals[3]+=Ma(Vals[6],Vals[4],Vals[5]);
+Vals[0]+=W[12];
+Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
+Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
+Vals[0]+=K[60];
+Vals[2]+=Vals[0];
+Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
+Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
 
 W[13]+=(rotr(W[14],7)^rotr(W[14],18)^(W[14]>>3U));
 W[13]+=W[6];
 W[13]+=(rotr(W[11],17)^rotr(W[11],19)^(W[11]>>10U));
-Vals[2]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
-Vals[2]+=ch(Vals[7],Vals[0],Vals[1]);
-Vals[2]+=K[61];
-Vals[2]+=W[13];
-Vals[6]+=Vals[2];
-Vals[2]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
-Vals[2]+=Ma(Vals[5],Vals[3],Vals[4]);
-
-W[14]+=(rotr(W[15],7)^rotr(W[15],18)^(W[15]>>3U));
-W[14]+=W[7];
-W[14]+=(rotr(W[12],17)^rotr(W[12],19)^(W[12]>>10U));
-Vals[1]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
-Vals[1]+=ch(Vals[6],Vals[7],Vals[0]);
-Vals[1]+=K[62];
-Vals[1]+=W[14];
-Vals[5]+=Vals[1];
-Vals[1]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
-Vals[1]+=Ma(Vals[4],Vals[2],Vals[3]);
-
-W[15]+=(rotr(W[0],7)^rotr(W[0],18)^(W[0]>>3U));
-W[15]+=W[8];
-W[15]+=(rotr(W[13],17)^rotr(W[13],19)^(W[13]>>10U));
-Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
-Vals[0]+=ch(Vals[5],Vals[6],Vals[7]);
-Vals[0]+=K[63];
-Vals[0]+=W[15];
-Vals[4]+=Vals[0];
-Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
-Vals[0]+=Ma(Vals[3],Vals[1],Vals[2]);
+Vals[6]+=W[13];
+Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
+Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
+Vals[6]+=K[61];
+Vals[3]+=Vals[6];
+Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
+Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
+
+Vals[7]+=W[14];
+Vals[7]+=(rotr(W[15],7)^rotr(W[15],18)^(W[15]>>3U));
+Vals[7]+=W[7];
+Vals[7]+=(rotr(W[12],17)^rotr(W[12],19)^(W[12]>>10U));
+Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
+Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
+Vals[7]+=K[62];
+Vals[4]+=Vals[7];
+Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
+Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
+
+Vals[5]+=W[15];
+Vals[5]+=(rotr(W[0],7)^rotr(W[0],18)^(W[0]>>3U));
+Vals[5]+=W[8];
+Vals[5]+=(rotr(W[13],17)^rotr(W[13],19)^(W[13]>>10U));
+Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
+Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
+Vals[5]+=K[63];
+Vals[1]+=Vals[5];
+Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
+Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
 
-W[0]=Vals[0];
-W[0]+=state0;
+Vals[5]+=state0;
 
 W[7]=state7;
-W[7]+=Vals[7];
+W[7]+=Vals[2];
 
-Vals[7]=0xF377ED68U;
-Vals[7]+=W[0];
+Vals[2]=0xF377ED68U;
+Vals[2]+=Vals[5];
 
 W[3]=state3;
-W[3]+=Vals[3];
+W[3]+=Vals[0];
 
-Vals[3]=0xa54ff53aU;
-Vals[3]+=Vals[7];
-Vals[7]+=0x08909ae5U;
+Vals[0]=0xa54ff53aU;
+Vals[0]+=Vals[2];
+Vals[2]+=0x08909ae5U;
 
 W[6]=state6;
-W[6]+=Vals[6];
+W[6]+=Vals[3];
 
-Vals[6]=0x90BB1E3CU;
-Vals[6]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
-Vals[6]+=(0x9b05688cU^(Vals[3]&0xca0b3af3U));
+Vals[3]=0x90BB1E3CU;
+Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
+Vals[3]+=(0x9b05688cU^(Vals[0]&0xca0b3af3U));
 
-W[1]=Vals[1];
-W[1]+=state1;
-Vals[6]+=W[1];
+Vals[7]+=state1;
+Vals[3]+=Vals[7];
 
 W[2]=state2;
-W[2]+=Vals[2];
+W[2]+=Vals[6];
 
-Vals[2]=0x3c6ef372U;
-Vals[2]+=Vals[6];
-Vals[6]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
-Vals[6]+=Ma2(0xbb67ae85U,Vals[7],0x6a09e667U);
+Vals[6]=0x3c6ef372U;
+Vals[6]+=Vals[3];
+Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
+Vals[3]+=Ma2(0xbb67ae85U,Vals[2],0x6a09e667U);
 
 W[5]=state5;
-W[5]+=Vals[5];
+W[5]+=Vals[4];
 
-Vals[5]=0x50C6645BU;
-Vals[5]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
-Vals[5]+=ch(Vals[2],Vals[3],0x510e527fU);
-Vals[5]+=W[2];
+Vals[4]=0x50C6645BU;
+Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
+Vals[4]+=ch(Vals[6],Vals[0],0x510e527fU);
+Vals[4]+=W[2];
 
-Vals[1]=0xbb67ae85U;
-Vals[1]+=Vals[5];
-Vals[5]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
-Vals[5]+=Ma2(0x6a09e667U,Vals[6],Vals[7]);
+W[1]=Vals[7];
+Vals[7]=0xbb67ae85U;
+Vals[7]+=Vals[4];
+Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
+Vals[4]+=Ma2(0x6a09e667U,Vals[3],Vals[2]);
 
 W[4]=state4;
-W[4]+=Vals[4];
+W[4]+=Vals[1];
 
-Vals[4]=0x3AC42E24U;
-Vals[4]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
-Vals[4]+=ch(Vals[1],Vals[2],Vals[3]);
-Vals[4]+=W[3];
+Vals[1]=0x3AC42E24U;
+Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
+Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
+Vals[1]+=W[3];
 
-Vals[0]=Vals[4];
-Vals[0]+=0x6a09e667U;
+W[0]=Vals[5];
 
-Vals[4]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
-Vals[4]+=Ma(Vals[7],Vals[5],Vals[6]);
+Vals[5]=Vals[1];
+Vals[5]+=0x6a09e667U;
 
-Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
-Vals[3]+=ch(Vals[0],Vals[1],Vals[2]);
-Vals[3]+=K[4];
-Vals[3]+=W[4];
-Vals[7]+=Vals[3];
-Vals[3]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
-Vals[3]+=Ma(Vals[6],Vals[4],Vals[5]);
-
-Vals[2]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
-Vals[2]+=ch(Vals[7],Vals[0],Vals[1]);
-Vals[2]+=K[5];
-Vals[2]+=W[5];
-Vals[6]+=Vals[2];
-Vals[2]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
-Vals[2]+=Ma(Vals[5],Vals[3],Vals[4]);
-
-Vals[1]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
-Vals[1]+=ch(Vals[6],Vals[7],Vals[0]);
-Vals[1]+=K[6];
-Vals[1]+=W[6];
-Vals[5]+=Vals[1];
-Vals[1]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
-Vals[1]+=Ma(Vals[4],Vals[2],Vals[3]);
+Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
+Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
 
 Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
-Vals[0]+=ch(Vals[5],Vals[6],Vals[7]);
-Vals[0]+=K[7];
-Vals[0]+=W[7];
-Vals[4]+=Vals[0];
+Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
+Vals[0]+=K[4];
+Vals[0]+=W[4];
+Vals[2]+=Vals[0];
 Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
-Vals[0]+=Ma(Vals[3],Vals[1],Vals[2]);
-
-Vals[7]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
-Vals[7]+=ch(Vals[4],Vals[5],Vals[6]);
-Vals[7]+=0x5807AA98U;
-Vals[3]+=Vals[7];
-Vals[7]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
-Vals[7]+=Ma(Vals[2],Vals[0],Vals[1]);
-
-Vals[6]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
-Vals[6]+=ch(Vals[3],Vals[4],Vals[5]);
-Vals[6]+=K[9];
-Vals[2]+=Vals[6];
-Vals[6]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
-Vals[6]+=Ma(Vals[1],Vals[7],Vals[0]);
-
-Vals[5]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
-Vals[5]+=ch(Vals[2],Vals[3],Vals[4]);
-Vals[5]+=K[10];
+Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
+
+Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
+Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
+Vals[6]+=K[5];
+Vals[6]+=W[5];
+Vals[3]+=Vals[6];
+Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
+Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
+
+Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
+Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
+Vals[7]+=K[6];
+Vals[7]+=W[6];
+Vals[4]+=Vals[7];
+Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
+Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
+
+Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
+Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
+Vals[5]+=K[7];
+Vals[5]+=W[7];
 Vals[1]+=Vals[5];
-Vals[5]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
-Vals[5]+=Ma(Vals[0],Vals[6],Vals[7]);
+Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
+Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
 
-Vals[4]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
-Vals[4]+=ch(Vals[1],Vals[2],Vals[3]);
-Vals[4]+=K[11];
-Vals[0]+=Vals[4];
-Vals[4]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
-Vals[4]+=Ma(Vals[7],Vals[5],Vals[6]);
+Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
+Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
+Vals[2]+=0x5807AA98U;
+Vals[0]+=Vals[2];
+Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
+Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
 
 Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
-Vals[3]+=ch(Vals[0],Vals[1],Vals[2]);
-Vals[3]+=K[12];
-Vals[7]+=Vals[3];
-Vals[3]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
-Vals[3]+=Ma(Vals[6],Vals[4],Vals[5]);
-
-Vals[2]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
-Vals[2]+=ch(Vals[7],Vals[0],Vals[1]);
-Vals[2]+=K[13];
-Vals[6]+=Vals[2];
-Vals[2]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
-Vals[2]+=Ma(Vals[5],Vals[3],Vals[4]);
-
-Vals[1]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
-Vals[1]+=ch(Vals[6],Vals[7],Vals[0]);
-Vals[1]+=K[14];
+Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
+Vals[3]+=K[9];
+Vals[6]+=Vals[3];
+Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
+Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
+
+Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
+Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
+Vals[4]+=K[10];
+Vals[7]+=Vals[4];
+Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
+Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
+
+Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
+Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
+Vals[1]+=K[11];
 Vals[5]+=Vals[1];
-Vals[1]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
-Vals[1]+=Ma(Vals[4],Vals[2],Vals[3]);
+Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
+Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
 
 Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
-Vals[0]+=ch(Vals[5],Vals[6],Vals[7]);
-Vals[0]+=0xC19BF274U;
-Vals[4]+=Vals[0];
+Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
+Vals[0]+=K[12];
+Vals[2]+=Vals[0];
 Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
-Vals[0]+=Ma(Vals[3],Vals[1],Vals[2]);
+Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
+
+Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
+Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
+Vals[6]+=K[13];
+Vals[3]+=Vals[6];
+Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
+Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
+
+Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
+Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
+Vals[7]+=K[14];
+Vals[4]+=Vals[7];
+Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
+Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
+
+Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
+Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
+Vals[5]+=0xC19BF274U;
+Vals[1]+=Vals[5];
+Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
+Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
 
 W[0]+=(rotr(W[1],7)^rotr(W[1],18)^(W[1]>>3U));
-Vals[7]+=W[0];
-Vals[7]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
-Vals[7]+=ch(Vals[4],Vals[5],Vals[6]);
-Vals[7]+=K[16];
-Vals[3]+=Vals[7];
-Vals[7]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
-Vals[7]+=Ma(Vals[2],Vals[0],Vals[1]);
+Vals[2]+=W[0];
+Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
+Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
+Vals[2]+=K[16];
+Vals[0]+=Vals[2];
+Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
+Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
 
 W[1]+=(rotr(W[2],7)^rotr(W[2],18)^(W[2]>>3U));
 W[1]+=0x00a00000U;
-Vals[6]+=W[1];
-Vals[6]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
-Vals[6]+=ch(Vals[3],Vals[4],Vals[5]);
-Vals[6]+=K[17];
-Vals[2]+=Vals[6];
-Vals[6]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
-Vals[6]+=Ma(Vals[1],Vals[7],Vals[0]);
+Vals[3]+=W[1];
+Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
+Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
+Vals[3]+=K[17];
+Vals[6]+=Vals[3];
+Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
+Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
 
 W[2]+=(rotr(W[3],7)^rotr(W[3],18)^(W[3]>>3U));
 W[2]+=(rotr(W[0],17)^rotr(W[0],19)^(W[0]>>10U));
-Vals[5]+=W[2];
-Vals[5]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
-Vals[5]+=ch(Vals[2],Vals[3],Vals[4]);
-Vals[5]+=K[18];
-Vals[1]+=Vals[5];
-Vals[5]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
-Vals[5]+=Ma(Vals[0],Vals[6],Vals[7]);
+Vals[4]+=W[2];
+Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
+Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
+Vals[4]+=K[18];
+Vals[7]+=Vals[4];
+Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
+Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
 
 W[3]+=(rotr(W[4],7)^rotr(W[4],18)^(W[4]>>3U));
 W[3]+=(rotr(W[1],17)^rotr(W[1],19)^(W[1]>>10U));
-Vals[4]+=W[3];
-Vals[4]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
-Vals[4]+=ch(Vals[1],Vals[2],Vals[3]);
-Vals[4]+=K[19];
-Vals[0]+=Vals[4];
-Vals[4]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
-Vals[4]+=Ma(Vals[7],Vals[5],Vals[6]);
+Vals[1]+=W[3];
+Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
+Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
+Vals[1]+=K[19];
+Vals[5]+=Vals[1];
+Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
+Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
 
 W[4]+=(rotr(W[5],7)^rotr(W[5],18)^(W[5]>>3U));
 W[4]+=(rotr(W[2],17)^rotr(W[2],19)^(W[2]>>10U));
-Vals[3]+=W[4];
-Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
-Vals[3]+=ch(Vals[0],Vals[1],Vals[2]);
-Vals[3]+=K[20];
-Vals[7]+=Vals[3];
-Vals[3]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
-Vals[3]+=Ma(Vals[6],Vals[4],Vals[5]);
+Vals[0]+=W[4];
+Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
+Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
+Vals[0]+=K[20];
+Vals[2]+=Vals[0];
+Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
+Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
 
 W[5]+=(rotr(W[6],7)^rotr(W[6],18)^(W[6]>>3U));
 W[5]+=(rotr(W[3],17)^rotr(W[3],19)^(W[3]>>10U));
-Vals[2]+=W[5];
-Vals[2]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
-Vals[2]+=ch(Vals[7],Vals[0],Vals[1]);
-Vals[2]+=K[21];
-Vals[6]+=Vals[2];
-Vals[2]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
-Vals[2]+=Ma(Vals[5],Vals[3],Vals[4]);
+Vals[6]+=W[5];
+Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
+Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
+Vals[6]+=K[21];
+Vals[3]+=Vals[6];
+Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
+Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
 
 W[6]+=(rotr(W[7],7)^rotr(W[7],18)^(W[7]>>3U));
 W[6]+=0x00000100U;
 W[6]+=(rotr(W[4],17)^rotr(W[4],19)^(W[4]>>10U));
-Vals[1]+=W[6];
-Vals[1]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
-Vals[1]+=ch(Vals[6],Vals[7],Vals[0]);
-Vals[1]+=K[22];
-Vals[5]+=Vals[1];
-Vals[1]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
-Vals[1]+=Ma(Vals[4],Vals[2],Vals[3]);
+Vals[7]+=W[6];
+Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
+Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
+Vals[7]+=K[22];
+Vals[4]+=Vals[7];
+Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
+Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
 
 W[7]+=0x11002000U;
 W[7]+=W[0];
 W[7]+=(rotr(W[5],17)^rotr(W[5],19)^(W[5]>>10U));
-Vals[0]+=W[7];
-Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
-Vals[0]+=ch(Vals[5],Vals[6],Vals[7]);
-Vals[0]+=K[23];
-Vals[4]+=Vals[0];
-Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
-Vals[0]+=Ma(Vals[3],Vals[1],Vals[2]);
+Vals[5]+=W[7];
+Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
+Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
+Vals[5]+=K[23];
+Vals[1]+=Vals[5];
+Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
+Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
 
 W[8]=0x80000000U;
 W[8]+=W[1];
 W[8]+=(rotr(W[6],17)^rotr(W[6],19)^(W[6]>>10U));
-Vals[7]+=W[8];
-Vals[7]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
-Vals[7]+=ch(Vals[4],Vals[5],Vals[6]);
-Vals[7]+=K[24];
-Vals[3]+=Vals[7];
-Vals[7]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
-Vals[7]+=Ma(Vals[2],Vals[0],Vals[1]);
+Vals[2]+=W[8];
+Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
+Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
+Vals[2]+=K[24];
+Vals[0]+=Vals[2];
+Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
+Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
 
 W[9]=W[2];
 W[9]+=(rotr(W[7],17)^rotr(W[7],19)^(W[7]>>10U));
-Vals[6]+=W[9];
-Vals[6]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
-Vals[6]+=ch(Vals[3],Vals[4],Vals[5]);
-Vals[6]+=K[25];
-Vals[2]+=Vals[6];
-Vals[6]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
-Vals[6]+=Ma(Vals[1],Vals[7],Vals[0]);
+Vals[3]+=W[9];
+Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
+Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
+Vals[3]+=K[25];
+Vals[6]+=Vals[3];
+Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
+Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
 
 W[10]=W[3];
 W[10]+=(rotr(W[8],17)^rotr(W[8],19)^(W[8]>>10U));
-Vals[5]+=W[10];
-Vals[5]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
-Vals[5]+=ch(Vals[2],Vals[3],Vals[4]);
-Vals[5]+=K[26];
-Vals[1]+=Vals[5];
-Vals[5]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
-Vals[5]+=Ma(Vals[0],Vals[6],Vals[7]);
+Vals[4]+=W[10];
+Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
+Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
+Vals[4]+=K[26];
+Vals[7]+=Vals[4];
+Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
+Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
 
 W[11]=W[4];
 W[11]+=(rotr(W[9],17)^rotr(W[9],19)^(W[9]>>10U));
-Vals[4]+=W[11];
-Vals[4]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
-Vals[4]+=ch(Vals[1],Vals[2],Vals[3]);
-Vals[4]+=K[27];
-Vals[0]+=Vals[4];
-Vals[4]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
-Vals[4]+=Ma(Vals[7],Vals[5],Vals[6]);
+Vals[1]+=W[11];
+Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
+Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
+Vals[1]+=K[27];
+Vals[5]+=Vals[1];
+Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
+Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
 
 W[12]=W[5];
 W[12]+=(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U));
-Vals[3]+=W[12];
-Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
-Vals[3]+=ch(Vals[0],Vals[1],Vals[2]);
-Vals[3]+=K[28];
-Vals[7]+=Vals[3];
-Vals[3]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
-Vals[3]+=Ma(Vals[6],Vals[4],Vals[5]);
+Vals[0]+=W[12];
+Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
+Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
+Vals[0]+=K[28];
+Vals[2]+=Vals[0];
+Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
+Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
 
 W[13]=W[6];
 W[13]+=(rotr(W[11],17)^rotr(W[11],19)^(W[11]>>10U));
-Vals[2]+=W[13];
-Vals[2]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
-Vals[2]+=ch(Vals[7],Vals[0],Vals[1]);
-Vals[2]+=K[29];
-Vals[6]+=Vals[2];
-Vals[2]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
-Vals[2]+=Ma(Vals[5],Vals[3],Vals[4]);
+Vals[6]+=W[13];
+Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
+Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
+Vals[6]+=K[29];
+Vals[3]+=Vals[6];
+Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
+Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
 
 W[14]=0x00400022U;
 W[14]+=W[7];
 W[14]+=(rotr(W[12],17)^rotr(W[12],19)^(W[12]>>10U));
-Vals[1]+=W[14];
-Vals[1]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
-Vals[1]+=ch(Vals[6],Vals[7],Vals[0]);
-Vals[1]+=K[30];
-Vals[5]+=Vals[1];
-Vals[1]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
-Vals[1]+=Ma(Vals[4],Vals[2],Vals[3]);
+Vals[7]+=W[14];
+Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
+Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
+Vals[7]+=K[30];
+Vals[4]+=Vals[7];
+Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
+Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
 
 W[15]=0x00000100U;
 W[15]+=(rotr(W[0],7)^rotr(W[0],18)^(W[0]>>3U));
 W[15]+=W[8];
 W[15]+=(rotr(W[13],17)^rotr(W[13],19)^(W[13]>>10U));
-Vals[0]+=W[15];
-Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
-Vals[0]+=ch(Vals[5],Vals[6],Vals[7]);
-Vals[0]+=K[31];
-Vals[4]+=Vals[0];
-Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
-Vals[0]+=Ma(Vals[3],Vals[1],Vals[2]);
+Vals[5]+=W[15];
+Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
+Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
+Vals[5]+=K[31];
+Vals[1]+=Vals[5];
+Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
+Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
 
 W[0]+=(rotr(W[1],7)^rotr(W[1],18)^(W[1]>>3U));
 W[0]+=W[9];
 W[0]+=(rotr(W[14],17)^rotr(W[14],19)^(W[14]>>10U));
-Vals[7]+=W[0];
-Vals[7]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
-Vals[7]+=ch(Vals[4],Vals[5],Vals[6]);
-Vals[7]+=K[32];
-Vals[3]+=Vals[7];
-Vals[7]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
-Vals[7]+=Ma(Vals[2],Vals[0],Vals[1]);
+Vals[2]+=W[0];
+Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
+Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
+Vals[2]+=K[32];
+Vals[0]+=Vals[2];
+Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
+Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
 
 W[1]+=(rotr(W[2],7)^rotr(W[2],18)^(W[2]>>3U));
 W[1]+=W[10];
 W[1]+=(rotr(W[15],17)^rotr(W[15],19)^(W[15]>>10U));
-Vals[6]+=W[1];
-Vals[6]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
-Vals[6]+=ch(Vals[3],Vals[4],Vals[5]);
-Vals[6]+=K[33];
-Vals[2]+=Vals[6];
-Vals[6]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
-Vals[6]+=Ma(Vals[1],Vals[7],Vals[0]);
+Vals[3]+=W[1];
+Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
+Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
+Vals[3]+=K[33];
+Vals[6]+=Vals[3];
+Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
+Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
 
 W[2]+=(rotr(W[3],7)^rotr(W[3],18)^(W[3]>>3U));
 W[2]+=W[11];
 W[2]+=(rotr(W[0],17)^rotr(W[0],19)^(W[0]>>10U));
-Vals[5]+=W[2];
-Vals[5]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
-Vals[5]+=ch(Vals[2],Vals[3],Vals[4]);
-Vals[5]+=K[34];
-Vals[1]+=Vals[5];
-Vals[5]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
-Vals[5]+=Ma(Vals[0],Vals[6],Vals[7]);
+Vals[4]+=W[2];
+Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
+Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
+Vals[4]+=K[34];
+Vals[7]+=Vals[4];
+Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
+Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
 
 W[3]+=(rotr(W[4],7)^rotr(W[4],18)^(W[4]>>3U));
 W[3]+=W[12];
 W[3]+=(rotr(W[1],17)^rotr(W[1],19)^(W[1]>>10U));
-Vals[4]+=W[3];
-Vals[4]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
-Vals[4]+=ch(Vals[1],Vals[2],Vals[3]);
-Vals[4]+=K[35];
-Vals[0]+=Vals[4];
-Vals[4]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
-Vals[4]+=Ma(Vals[7],Vals[5],Vals[6]);
+Vals[1]+=W[3];
+Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
+Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
+Vals[1]+=K[35];
+Vals[5]+=Vals[1];
+Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
+Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
 
 W[4]+=(rotr(W[5],7)^rotr(W[5],18)^(W[5]>>3U));
 W[4]+=W[13];
 W[4]+=(rotr(W[2],17)^rotr(W[2],19)^(W[2]>>10U));
-Vals[3]+=W[4];
-Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
-Vals[3]+=ch(Vals[0],Vals[1],Vals[2]);
-Vals[3]+=K[36];
-Vals[7]+=Vals[3];
-Vals[3]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
-Vals[3]+=Ma(Vals[6],Vals[4],Vals[5]);
+Vals[0]+=W[4];
+Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
+Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
+Vals[0]+=K[36];
+Vals[2]+=Vals[0];
+Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
+Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
 
 W[5]+=(rotr(W[6],7)^rotr(W[6],18)^(W[6]>>3U));
 W[5]+=W[14];
 W[5]+=(rotr(W[3],17)^rotr(W[3],19)^(W[3]>>10U));
-Vals[2]+=W[5];
-Vals[2]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
-Vals[2]+=ch(Vals[7],Vals[0],Vals[1]);
-Vals[2]+=K[37];
-Vals[6]+=Vals[2];
-Vals[2]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
-Vals[2]+=Ma(Vals[5],Vals[3],Vals[4]);
+Vals[6]+=W[5];
+Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
+Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
+Vals[6]+=K[37];
+Vals[3]+=Vals[6];
+Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
+Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
 
 W[6]+=(rotr(W[7],7)^rotr(W[7],18)^(W[7]>>3U));
 W[6]+=W[15];
 W[6]+=(rotr(W[4],17)^rotr(W[4],19)^(W[4]>>10U));
-Vals[1]+=W[6];
-Vals[1]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
-Vals[1]+=ch(Vals[6],Vals[7],Vals[0]);
-Vals[1]+=K[38];
-Vals[5]+=Vals[1];
-Vals[1]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
-Vals[1]+=Ma(Vals[4],Vals[2],Vals[3]);
+Vals[7]+=W[6];
+Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
+Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
+Vals[7]+=K[38];
+Vals[4]+=Vals[7];
+Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
+Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
 
 W[7]+=(rotr(W[8],7)^rotr(W[8],18)^(W[8]>>3U));
 W[7]+=W[0];
 W[7]+=(rotr(W[5],17)^rotr(W[5],19)^(W[5]>>10U));
-Vals[0]+=W[7];
-Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
-Vals[0]+=ch(Vals[5],Vals[6],Vals[7]);
-Vals[0]+=K[39];
-Vals[4]+=Vals[0];
-Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
-Vals[0]+=Ma(Vals[3],Vals[1],Vals[2]);
+Vals[5]+=W[7];
+Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
+Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
+Vals[5]+=K[39];
+Vals[1]+=Vals[5];
+Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
+Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
 
 W[8]+=(rotr(W[9],7)^rotr(W[9],18)^(W[9]>>3U));
 W[8]+=W[1];
 W[8]+=(rotr(W[6],17)^rotr(W[6],19)^(W[6]>>10U));
-Vals[7]+=W[8];
-Vals[7]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
-Vals[7]+=ch(Vals[4],Vals[5],Vals[6]);
-Vals[7]+=K[40];
-Vals[3]+=Vals[7];
-Vals[7]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
-Vals[7]+=Ma(Vals[2],Vals[0],Vals[1]);
+Vals[2]+=W[8];
+Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
+Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
+Vals[2]+=K[40];
+Vals[0]+=Vals[2];
+Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
+Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
 
 W[9]+=(rotr(W[10],7)^rotr(W[10],18)^(W[10]>>3U));
 W[9]+=W[2];
 W[9]+=(rotr(W[7],17)^rotr(W[7],19)^(W[7]>>10U));
-Vals[6]+=W[9];
-Vals[6]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
-Vals[6]+=ch(Vals[3],Vals[4],Vals[5]);
-Vals[6]+=K[41];
-Vals[2]+=Vals[6];
-Vals[6]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
-Vals[6]+=Ma(Vals[1],Vals[7],Vals[0]);
+Vals[3]+=W[9];
+Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
+Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
+Vals[3]+=K[41];
+Vals[6]+=Vals[3];
+Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
+Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
 
 W[10]+=(rotr(W[11],7)^rotr(W[11],18)^(W[11]>>3U));
 W[10]+=W[3];
 W[10]+=(rotr(W[8],17)^rotr(W[8],19)^(W[8]>>10U));
-Vals[5]+=W[10];
-Vals[5]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
-Vals[5]+=ch(Vals[2],Vals[3],Vals[4]);
-Vals[5]+=K[42];
-Vals[1]+=Vals[5];
-Vals[5]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
-Vals[5]+=Ma(Vals[0],Vals[6],Vals[7]);
+Vals[4]+=W[10];
+Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
+Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
+Vals[4]+=K[42];
+Vals[7]+=Vals[4];
+Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
+Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
 
 W[11]+=(rotr(W[12],7)^rotr(W[12],18)^(W[12]>>3U));
 W[11]+=W[4];
 W[11]+=(rotr(W[9],17)^rotr(W[9],19)^(W[9]>>10U));
-Vals[4]+=W[11];
-Vals[4]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
-Vals[4]+=ch(Vals[1],Vals[2],Vals[3]);
-Vals[4]+=K[43];
-Vals[0]+=Vals[4];
-Vals[4]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
-Vals[4]+=Ma(Vals[7],Vals[5],Vals[6]);
+Vals[1]+=W[11];
+Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
+Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
+Vals[1]+=K[43];
+Vals[5]+=Vals[1];
+Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
+Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
 
 W[12]+=(rotr(W[13],7)^rotr(W[13],18)^(W[13]>>3U));
 W[12]+=W[5];
 W[12]+=(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U));
-Vals[3]+=W[12];
-Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
-Vals[3]+=ch(Vals[0],Vals[1],Vals[2]);
-Vals[3]+=K[44];
-Vals[7]+=Vals[3];
-Vals[3]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
-Vals[3]+=Ma(Vals[6],Vals[4],Vals[5]);
+Vals[0]+=W[12];
+Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
+Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
+Vals[0]+=K[44];
+Vals[2]+=Vals[0];
+Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
+Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
 
 W[13]+=(rotr(W[14],7)^rotr(W[14],18)^(W[14]>>3U));
 W[13]+=W[6];
 W[13]+=(rotr(W[11],17)^rotr(W[11],19)^(W[11]>>10U));
-Vals[2]+=W[13];
-Vals[2]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
-Vals[2]+=ch(Vals[7],Vals[0],Vals[1]);
-Vals[2]+=K[45];
-Vals[6]+=Vals[2];
-Vals[2]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
-Vals[2]+=Ma(Vals[5],Vals[3],Vals[4]);
+Vals[6]+=W[13];
+Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
+Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
+Vals[6]+=K[45];
+Vals[3]+=Vals[6];
+Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
+Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
 
 W[14]+=(rotr(W[15],7)^rotr(W[15],18)^(W[15]>>3U));
 W[14]+=W[7];
 W[14]+=(rotr(W[12],17)^rotr(W[12],19)^(W[12]>>10U));
-Vals[1]+=W[14];
-Vals[1]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
-Vals[1]+=ch(Vals[6],Vals[7],Vals[0]);
-Vals[1]+=K[46];
-Vals[5]+=Vals[1];
-Vals[1]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
-Vals[1]+=Ma(Vals[4],Vals[2],Vals[3]);
+Vals[7]+=W[14];
+Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
+Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
+Vals[7]+=K[46];
+Vals[4]+=Vals[7];
+Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
+Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
 
 W[15]+=(rotr(W[0],7)^rotr(W[0],18)^(W[0]>>3U));
 W[15]+=W[8];
 W[15]+=(rotr(W[13],17)^rotr(W[13],19)^(W[13]>>10U));
-Vals[0]+=W[15];
-Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
-Vals[0]+=ch(Vals[5],Vals[6],Vals[7]);
-Vals[0]+=K[47];
-Vals[4]+=Vals[0];
-Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
-Vals[0]+=Ma(Vals[3],Vals[1],Vals[2]);
+Vals[5]+=W[15];
+Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
+Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
+Vals[5]+=K[47];
+Vals[1]+=Vals[5];
+Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
+Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
 
 W[0]+=(rotr(W[1],7)^rotr(W[1],18)^(W[1]>>3U));
 W[0]+=W[9];
 W[0]+=(rotr(W[14],17)^rotr(W[14],19)^(W[14]>>10U));
-Vals[7]+=W[0];
-Vals[7]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
-Vals[7]+=ch(Vals[4],Vals[5],Vals[6]);
-Vals[7]+=K[48];
-Vals[3]+=Vals[7];
-Vals[7]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
-Vals[7]+=Ma(Vals[2],Vals[0],Vals[1]);
+Vals[2]+=W[0];
+Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
+Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
+Vals[2]+=K[48];
+Vals[0]+=Vals[2];
+Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
+Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
 
 W[1]+=(rotr(W[2],7)^rotr(W[2],18)^(W[2]>>3U));
 W[1]+=W[10];
 W[1]+=(rotr(W[15],17)^rotr(W[15],19)^(W[15]>>10U));
-Vals[6]+=W[1];
-Vals[6]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
-Vals[6]+=ch(Vals[3],Vals[4],Vals[5]);
-Vals[6]+=K[49];
-Vals[2]+=Vals[6];
-Vals[6]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
-Vals[6]+=Ma(Vals[1],Vals[7],Vals[0]);
+Vals[3]+=W[1];
+Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
+Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
+Vals[3]+=K[49];
+Vals[6]+=Vals[3];
+Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
+Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
 
 W[2]+=(rotr(W[3],7)^rotr(W[3],18)^(W[3]>>3U));
 W[2]+=W[11];
 W[2]+=(rotr(W[0],17)^rotr(W[0],19)^(W[0]>>10U));
-Vals[5]+=W[2];
-Vals[5]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
-Vals[5]+=ch(Vals[2],Vals[3],Vals[4]);
-Vals[5]+=K[50];
-Vals[1]+=Vals[5];
-Vals[5]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
-Vals[5]+=Ma(Vals[0],Vals[6],Vals[7]);
+Vals[4]+=W[2];
+Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
+Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
+Vals[4]+=K[50];
+Vals[7]+=Vals[4];
+Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
+Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
 
 W[3]+=(rotr(W[4],7)^rotr(W[4],18)^(W[4]>>3U));
 W[3]+=W[12];
 W[3]+=(rotr(W[1],17)^rotr(W[1],19)^(W[1]>>10U));
-Vals[4]+=W[3];
-Vals[4]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
-Vals[4]+=ch(Vals[1],Vals[2],Vals[3]);
-Vals[4]+=K[51];
-Vals[0]+=Vals[4];
-Vals[4]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
-Vals[4]+=Ma(Vals[7],Vals[5],Vals[6]);
+Vals[1]+=W[3];
+Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
+Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
+Vals[1]+=K[51];
+Vals[5]+=Vals[1];
+Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
+Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
 
 W[4]+=(rotr(W[5],7)^rotr(W[5],18)^(W[5]>>3U));
 W[4]+=W[13];
 W[4]+=(rotr(W[2],17)^rotr(W[2],19)^(W[2]>>10U));
-Vals[3]+=W[4];
-Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
-Vals[3]+=ch(Vals[0],Vals[1],Vals[2]);
-Vals[3]+=K[52];
-Vals[7]+=Vals[3];
-Vals[3]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
-Vals[3]+=Ma(Vals[6],Vals[4],Vals[5]);
+Vals[0]+=W[4];
+Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
+Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
+Vals[0]+=K[52];
+Vals[2]+=Vals[0];
+Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
+Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
 
 W[5]+=(rotr(W[6],7)^rotr(W[6],18)^(W[6]>>3U));
 W[5]+=W[14];
 W[5]+=(rotr(W[3],17)^rotr(W[3],19)^(W[3]>>10U));
-Vals[2]+=W[5];
-Vals[2]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
-Vals[2]+=ch(Vals[7],Vals[0],Vals[1]);
-Vals[2]+=K[53];
-Vals[6]+=Vals[2];
-Vals[2]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
-Vals[2]+=Ma(Vals[5],Vals[3],Vals[4]);
+Vals[6]+=W[5];
+Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
+Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
+Vals[6]+=K[53];
+Vals[3]+=Vals[6];
+Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
+Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
 
 W[6]+=(rotr(W[7],7)^rotr(W[7],18)^(W[7]>>3U));
 W[6]+=W[15];
 W[6]+=(rotr(W[4],17)^rotr(W[4],19)^(W[4]>>10U));
-Vals[1]+=W[6];
-Vals[1]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
-Vals[1]+=ch(Vals[6],Vals[7],Vals[0]);
-Vals[1]+=K[54];
-Vals[5]+=Vals[1];
-Vals[1]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
-Vals[1]+=Ma(Vals[4],Vals[2],Vals[3]);
+Vals[7]+=W[6];
+Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
+Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
+Vals[7]+=K[54];
+Vals[4]+=Vals[7];
+Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
+Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
 
 W[7]+=(rotr(W[8],7)^rotr(W[8],18)^(W[8]>>3U));
 W[7]+=W[0];
 W[7]+=(rotr(W[5],17)^rotr(W[5],19)^(W[5]>>10U));
-Vals[0]+=W[7];
-Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
-Vals[0]+=ch(Vals[5],Vals[6],Vals[7]);
-Vals[0]+=K[55];
-Vals[4]+=Vals[0];
-Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
-Vals[0]+=Ma(Vals[3],Vals[1],Vals[2]);
+Vals[5]+=W[7];
+Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
+Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
+Vals[5]+=K[55];
+Vals[1]+=Vals[5];
+Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
+Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
 
 W[8]+=(rotr(W[9],7)^rotr(W[9],18)^(W[9]>>3U));
 W[8]+=W[1];
 W[8]+=(rotr(W[6],17)^rotr(W[6],19)^(W[6]>>10U));
-Vals[7]+=W[8];
-Vals[7]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
-Vals[7]+=ch(Vals[4],Vals[5],Vals[6]);
-Vals[7]+=K[56];
-Vals[3]+=Vals[7];
+Vals[2]+=W[8];
+Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
+Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
+Vals[2]+=K[56];
+Vals[0]+=Vals[2];
 
 W[9]+=(rotr(W[10],7)^rotr(W[10],18)^(W[10]>>3U));
 W[9]+=W[2];
 W[9]+=(rotr(W[7],17)^rotr(W[7],19)^(W[7]>>10U));
-Vals[6]+=W[9];
-Vals[6]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
-Vals[6]+=ch(Vals[3],Vals[4],Vals[5]);
-Vals[6]+=K[57];
-Vals[6]+=Vals[2];
+Vals[3]+=W[9];
+Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
+Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
+Vals[3]+=K[57];
+Vals[3]+=Vals[6];
 
 W[10]+=(rotr(W[11],7)^rotr(W[11],18)^(W[11]>>3U));
 W[10]+=W[3];
 W[10]+=(rotr(W[8],17)^rotr(W[8],19)^(W[8]>>10U));
-Vals[5]+=W[10];
-Vals[5]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
-Vals[5]+=ch(Vals[6],Vals[3],Vals[4]);
-Vals[5]+=K[58];
-Vals[5]+=Vals[1];
-Vals[4]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
-Vals[4]+=ch(Vals[5],Vals[6],Vals[3]);
-Vals[4]+=W[11];
-Vals[4]+=(rotr(W[12],7)^rotr(W[12],18)^(W[12]>>3U));
-Vals[4]+=W[4];
-Vals[4]+=(rotr(W[9],17)^rotr(W[9],19)^(W[9]>>10U));
-Vals[4]+=K[59];
-Vals[4]+=Vals[0];
+Vals[4]+=W[10];
+Vals[4]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
+Vals[4]+=ch(Vals[3],Vals[0],Vals[1]);
+Vals[4]+=K[58];
+Vals[4]+=Vals[7];
+Vals[1]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
+Vals[1]+=ch(Vals[4],Vals[3],Vals[0]);
+Vals[1]+=W[11];
+Vals[1]+=(rotr(W[12],7)^rotr(W[12],18)^(W[12]>>3U));
+Vals[1]+=W[4];
+Vals[1]+=(rotr(W[9],17)^rotr(W[9],19)^(W[9]>>10U));
+Vals[1]+=K[59];
+Vals[1]+=Vals[5];
 
 #define FOUND (0x80)
 #define NFLAG (0x7F)
 
 #if defined(VECTORS2) || defined(VECTORS4)
-	Vals[7]+=Ma(Vals[2],Vals[0],Vals[1]);
-	Vals[7]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
-	Vals[7]+=W[12];
-	Vals[7]+=(rotr(W[13],7)^rotr(W[13],18)^(W[13]>>3U));
-	Vals[7]+=W[5];
-	Vals[7]+=(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U));
-	Vals[7]+=Vals[3];
-	Vals[7]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
-	Vals[7]+=ch(Vals[4],Vals[5],Vals[6]);
-
-	if (any(Vals[7] == 0x136032edU)) {
-		if (Vals[7].x == 0x136032edU)
+	Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
+	Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
+	Vals[2]+=W[12];
+	Vals[2]+=(rotr(W[13],7)^rotr(W[13],18)^(W[13]>>3U));
+	Vals[2]+=W[5];
+	Vals[2]+=(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U));
+	Vals[2]+=Vals[0];
+	Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
+	Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
+
+	if (any(Vals[2] == 0x136032edU)) {
+		if (Vals[2].x == 0x136032edU)
 			output[FOUND] = output[NFLAG & nonce.x] = nonce.x;
-		if (Vals[7].y == 0x136032edU)
+		if (Vals[2].y == 0x136032edU)
 			output[FOUND] = output[NFLAG & nonce.y] = nonce.y;
 #if defined(VECTORS4)
-		if (Vals[7].z == 0x136032edU)
+		if (Vals[2].z == 0x136032edU)
 			output[FOUND] = output[NFLAG & nonce.z] = nonce.z;
-		if (Vals[7].w == 0x136032edU)
+		if (Vals[2].w == 0x136032edU)
 			output[FOUND] = output[NFLAG & nonce.w] = nonce.w;
 #endif
 	}
 #else
-	if ((Vals[7]+
-		Ma(Vals[2],Vals[0],Vals[1])+
-		(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22))+
+	if ((Vals[2]+
+		Ma(Vals[6],Vals[5],Vals[7])+
+		(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22))+
 		W[12]+
 		(rotr(W[13],7)^rotr(W[13],18)^(W[13]>>3U))+
 		W[5]+
 		(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U))+
-		Vals[3]+
-		(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25))+
-		ch(Vals[4],Vals[5],Vals[6])) == 0x136032edU)
+		Vals[0]+
+		(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25))+
+		ch(Vals[1],Vals[4],Vals[3])) == 0x136032edU)
 			output[FOUND] = output[NFLAG & nonce] =  nonce;
 #endif
 }

sha256_4way.c

@@ -4,7 +4,7 @@
 
 // tcatm's 4-way 128-bit SSE2 SHA-256
 
-#include "device-cpu.h"
+#include "driver-cpu.h"
 
 #ifdef WANT_SSE2_4WAY
 
@@ -111,8 +111,6 @@ bool ScanHash_4WaySSE2(int thr_id, const unsigned char *pmidstate,
 
 	pdata += 64;
 
-    work_restart[thr_id].restart = 0;
-
     for (;;)
     {
         unsigned int thash[9][NPAR] __attribute__((aligned(128)));

sha256_altivec_4way.c

@@ -9,7 +9,7 @@
 //
 
 
-#include "device-cpu.h"
+#include "driver-cpu.h"
 
 #ifdef WANT_ALTIVEC_4WAY
 
@@ -84,8 +84,6 @@ bool ScanHash_altivec_4way(int thr_id, const unsigned char *pmidstate,
 
 	pdata += 64;
 
-    work_restart[thr_id].restart = 0;
-
     for (;;)
     {
         unsigned int thash[9][NPAR] __attribute__((aligned(128)));

sha256_cryptopp.c

@@ -589,8 +589,6 @@ bool scanhash_asm32(int thr_id, const unsigned char *midstate,
 
 	data += 64;
 
-	work_restart[thr_id].restart = 0;
-
 	while (1) {
 		n++;
 		*nonce = n;

sha256_generic.c

@@ -251,8 +251,6 @@ bool scanhash_c(int thr_id, const unsigned char *midstate, unsigned char *data,
 
 	data += 64;
 
-	work_restart[thr_id].restart = 0;
-
 	while (1) {
 		n++;
 		*nonce = n;

sha256_sse2_amd64.c

@@ -9,7 +9,7 @@
  *
  */
 
-#include "device-cpu.h"
+#include "driver-cpu.h"
 
 #ifdef WANT_X8664_SSE2
 
@@ -65,8 +65,6 @@ bool scanhash_sse2_64(int thr_id, const unsigned char *pmidstate,
 
 	pdata += 64;
 
-    work_restart[thr_id].restart = 0;
-
     /* For debugging */
     union {
         __m128i m;

sha256_sse2_i386.c

@@ -9,7 +9,7 @@
  *
  */
 
-#include "device-cpu.h"
+#include "driver-cpu.h"
 
 #ifdef WANT_X8632_SSE2
 
@@ -65,8 +65,6 @@ bool scanhash_sse2_32(int thr_id, const unsigned char *pmidstate,
 
 	pdata += 64;
 
-    work_restart[thr_id].restart = 0;
-
     /* Message expansion */
     memcpy(m_midstate, pmidstate, sizeof(m_midstate));
     memcpy(m_w, pdata, sizeof(m_w)); /* The 2nd half of the data */

sha256_sse4_amd64.c

@@ -9,7 +9,7 @@
  *
  */
 
-#include "device-cpu.h"
+#include "driver-cpu.h"
 
 #ifdef WANT_X8664_SSE4
 
@@ -62,8 +62,6 @@ bool scanhash_sse4_64(int thr_id, const unsigned char *pmidstate,
 
 	pdata += 64;
 
-    work_restart[thr_id].restart = 0;
-
     /* For debugging */
     union {
         __m128i m;

sha256_via.c

@@ -1,5 +1,5 @@
 
-#include "device-cpu.h"
+#include "driver-cpu.h"
 
 #include <stdint.h>
 #include <stdlib.h>
@@ -35,8 +35,6 @@ bool scanhash_via(int thr_id, const unsigned char *pmidstate,
 	unsigned long stat_ctr = 0;
 	int i;
 
-	work_restart[thr_id].restart = 0;
-
 	/* bitcoin gives us big endian input, but via wants LE,
 	 * so we reverse the swapping bitcoin has already done (extra work)
 	 * in order to permit the hardware to swap everything

todo_ztex.txt

@@ -0,0 +1,5 @@
+- verify setting cgpu.status=DEAD does in fact stop the thread
+- allow configuring bitstream directory
+- HS fpga config
+- allow configuring LIBZTEX_OVERHEATTHRESHOLD
+- hotplug support?

util.c

@@ -30,6 +30,7 @@
 # include <winsock2.h>
 # include <mstcpip.h>
 #endif
+
 #include "miner.h"
 #include "elist.h"
 #include "compat.h"
@@ -364,10 +365,16 @@ json_t *json_rpc_call(CURL *curl, const char *url,
 	if (probing) {
 		pool->probed = true;
 		/* If X-Long-Polling was found, activate long polling */
-		if (hi.lp_path)
+		if (hi.lp_path) {
+			if (pool->hdr_path != NULL)
+				free(pool->hdr_path);
 			pool->hdr_path = hi.lp_path;
-		else
+		} else {
 			pool->hdr_path = NULL;
+		}
+	} else if (hi.lp_path) {
+		free(hi.lp_path);
+		hi.lp_path = NULL;
 	}
 
 	*rolltime = hi.has_rolltime;
@@ -410,9 +417,11 @@ json_t *json_rpc_call(CURL *curl, const char *url,
 		goto err_out;
 	}
 
-	if (hi.reason)
+	if (hi.reason) {
 		json_object_set_new(val, "reject-reason", json_string(hi.reason));
-
+		free(hi.reason);
+		hi.reason = NULL;
+	}
 	successful_connect = true;
 	databuf_free(&all_data);
 	curl_slist_free_all(headers);