jgarzik/cpuminer — pull request #56 by MasterCATZ (2014-02-03 08:02:33+00:00)
Title: add scypt-jane support to 3.7.2

diff -urN cgminer-3.7.2.orig/cgminer.c cgminer-3.7.2/cgminer.c
--- cgminer-3.7.2.orig/cgminer.c	2013-11-06 00:24:07.000000000 +1100
+++ cgminer-3.7.2/cgminer.c	2014-01-02 17:30:53.821237874 +1100
@@ -58,6 +58,7 @@
#include "driver-opencl.h"
#include "bench_block.h"
#include "scrypt.h"
+#include "scrypt-jane.h"
#ifdef USE_USBUTILS
#include "usbutils.h"
#endif
@@ -113,6 +114,11 @@
unsigned long long global_hashrate;
unsigned long global_quota_gcd = 1;

+/* scrypt-jane, defaults suitable for YAC */
+unsigned int sj_minNf = 4;
+unsigned int sj_maxNf = 30;
+unsigned int sj_startTime = 1388361600;
+
#if defined(HAVE_OPENCL) || defined(USE_USBUTILS)
int nDevs;
#endif
@@ -122,6 +128,7 @@
int gpu_threads;
#ifdef USE_SCRYPT
bool opt_scrypt;
+bool opt_scrypt_jane;
#endif
#endif
bool opt_restart = true;
@@ -650,6 +657,11 @@
return set_int_range(arg, i, 0, 10);
}

+static char *set_int_0_to_40(const char *arg, int *i)
+{
+	return set_int_range(arg, i, 0, 40);
+}
+
#ifdef USE_AVALON
static char *set_int_0_to_100(const char *arg, int *i)
{
@@ -1373,6 +1385,18 @@
OPT_WITHOUT_ARG("--scrypt",
opt_set_bool, &opt_scrypt,
"Use the scrypt algorithm for mining (litecoin only)"),
+	OPT_WITHOUT_ARG("--scrypt-jane",
+	opt_set_bool, &opt_scrypt_jane,
+	"Use the scrypt-jane algorithm for mining"),
+	OPT_WITH_ARG("--sj-nfmin",
+	set_int_0_to_40, opt_show_intval, &sj_minNf,
+	"Set min N factor for mining scrypt-jane"),
+	OPT_WITH_ARG("--sj-nfmax",
+	set_int_0_to_40, opt_show_intval, &sj_maxNf,
+	"Set max N factor for mining scrypt-jane"),
+	OPT_WITH_ARG("--sj-time",
+	opt_set_intval, opt_show_intval, &sj_startTime,
+	"Set StartTime for mining scrypt-jane"),
OPT_WITH_ARG("--shaders",
set_shaders, NULL, NULL,
"GPU shaders per card for tuning scrypt, comma separated"),
@@ -3802,10 +3826,17 @@

static void rebuild_hash(struct work *work)
{
-	if (opt_scrypt)
-	scrypt_regenhash(work);
-	else
+	if (opt_scrypt) {
+	if (opt_scrypt_jane) {
+	/* scrypt-jane */
+	sj_scrypt_regenhash(work);
+	} else {
+	/* standard scrypt */
+	scrypt_regenhash(work);
+	}
+	} else {
regen_hash(work);
+	}
}

static bool cnx_needed(struct pool *pool);
@@ -4480,6 +4511,9 @@
case KL_SCRYPT:
fprintf(fcfg, "scrypt");
break;
+	case KL_SCRYPT_JANE:
+	fprintf(fcfg, "scrypt-jane");
+	break;
}
}
#ifdef USE_SCRYPT
@@ -6202,7 +6236,11 @@
{
uint32_t *work_nonce = (uint32_t *)(work->data + 64 + 12);

-	*work_nonce = htole32(nonce);
+	if (opt_scrypt_jane) {
+	*work_nonce = htobe32(nonce);
+	} else {
+	*work_nonce = htole32(nonce);
+	}

rebuild_hash(work);
}
@@ -8114,6 +8152,10 @@
if (!config_loaded)
load_default_config();

+	if (opt_scrypt_jane)
+	opt_set_bool(&opt_scrypt);
+
+
if (opt_benchmark) {
struct pool *pool;

diff -urN cgminer-3.7.2.orig/config.h.in cgminer-3.7.2/config.h.in
--- cgminer-3.7.2.orig/config.h.in	2013-11-06 00:29:24.000000000 +1100
+++ cgminer-3.7.2/config.h.in	2014-01-02 17:30:53.829238126 +1100
@@ -334,6 +334,9 @@
/* Filename for scrypt kernel */
#undef SCRYPT_KERNNAME

+/* Filename for scrypt-jane kernel */
+#undef SCRYPT_JANE_KERNNAME
+
/* Define to l, ll, u, ul, ull, etc., as suitable for constants of type
'sig_atomic_t'. */
#undef SIG_ATOMIC_T_SUFFIX
diff -urN cgminer-3.7.2.orig/configure cgminer-3.7.2/configure
--- cgminer-3.7.2.orig/configure	2013-11-06 00:29:23.000000000 +1100
+++ cgminer-3.7.2/configure	2014-01-02 17:30:53.881239099 +1100
@@ -17333,6 +17333,11 @@
_ACEOF

+cat >>confdefs.h <<_ACEOF
+#define SCRYPT_JANE_KERNNAME "scrypt-jane"
+_ACEOF
+
+

diff -urN cgminer-3.7.2.orig/configure.ac cgminer-3.7.2/configure.ac
--- cgminer-3.7.2.orig/configure.ac	2013-11-06 00:26:04.000000000 +1100
+++ cgminer-3.7.2/configure.ac	2014-01-02 17:30:53.889239262 +1100
@@ -481,6 +481,7 @@
AC_DEFINE_UNQUOTED([DIAKGCN_KERNNAME], ["diakgcn121016"], [Filename for diakgcn kernel])
AC_DEFINE_UNQUOTED([DIABLO_KERNNAME], ["diablo130302"], [Filename for diablo kernel])
AC_DEFINE_UNQUOTED([SCRYPT_KERNNAME], ["scrypt130511"], [Filename for scrypt kernel])
+AC_DEFINE_UNQUOTED([SCRYPT_JANE_KERNNAME], ["scrypt-jane"], [Filename for scrypt-jane kernel])

AC_SUBST(OPENCL_LIBS)
diff -urN cgminer-3.7.2.orig/driver-opencl.c cgminer-3.7.2/driver-opencl.c
--- cgminer-3.7.2.orig/driver-opencl.c	2013-10-08 23:35:01.000000000 +1100
+++ cgminer-3.7.2/driver-opencl.c	2014-01-02 17:30:53.897239423 +1100
@@ -34,6 +34,10 @@
#include "adl.h"
#include "util.h"

+#ifdef USE_SCRYPT
+#include "scrypt-jane.h"	/* sj_be32enc_vect */
+#endif
+
/* TODO: cleanup externals ********************/

#ifdef HAVE_CURSES
@@ -211,6 +215,8 @@
#ifdef USE_SCRYPT
if (!strcmp(arg, "scrypt"))
return KL_SCRYPT;
+	if (!strcmp(arg, "scrypt-jane"))
+	return KL_SCRYPT_JANE;
#endif
return KL_NONE;
}
@@ -1067,6 +1073,7 @@
}

#ifdef USE_SCRYPT
+
static cl_int queue_scrypt_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_unused cl_uint threads)
{
unsigned char *midstate = blk->work->midstate;
@@ -1074,9 +1081,26 @@
unsigned int num = 0;
cl_uint le_target;
cl_int status = 0;
+	uint32_t data[20];
+	unsigned int timestamp;
+	cl_uint nfactor;
+
+	if (opt_scrypt_jane) {
+	timestamp = bswap_32(*((unsigned int *)(blk->work->data + 17*4)));
+	nfactor = sj_GetNfactor(timestamp);
+	nfactor = (1 << (nfactor + 1));
+	}

le_target = *(cl_uint *)(blk->work->device_target + 28);
-	clState->cldata = blk->work->data;
+	if (!opt_scrypt_jane) {
+	clState->cldata = blk->work->data;
+	} else {
+	/* scrypt-jane */
+	applog(LOG_DEBUG, "Timestamp: %d, Nfactor: %d, Target: %x", timestamp, nfactor, le_target);
+	sj_be32enc_vect(data, (const uint32_t *)blk->work->data, 19);
+	clState->cldata = data;
+	}
+
status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 80, clState->cldata, 0, NULL,NULL);

CL_SET_ARG(clState->CLbuffer0);
@@ -1085,6 +1109,8 @@
CL_SET_VARG(4, &midstate[0]);
CL_SET_VARG(4, &midstate[16]);
CL_SET_ARG(le_target);
+	if (opt_scrypt_jane)
+	CL_SET_ARG(nfactor);

return status;
}
@@ -1378,6 +1404,9 @@
case KL_SCRYPT:
cgpu->kname = "scrypt";
break;
+	case KL_SCRYPT_JANE:
+	cgpu->kname = "scrypt-jane";
+	break;
#endif
case KL_POCLBM:
cgpu->kname = "poclbm";
@@ -1423,6 +1452,7 @@
break;
#ifdef USE_SCRYPT
case KL_SCRYPT:
+	case KL_SCRYPT_JANE:
thrdata->queue_kernel_parameters = &queue_scrypt_kernel;
break;
#endif
@@ -1484,6 +1514,8 @@
int64_t hashes;
int found = opt_scrypt ? SCRYPT_FOUND : FOUND;
int buffersize = opt_scrypt ? SCRYPT_BUFFERSIZE : BUFFERSIZE;
+	uint32_t *o;
+	uint32_t target;

/* Windows' timer resolution is only 15ms so oversample 5x */
if (gpu->dynamic && (++gpu->intervals * dynamic_us) > 70000) {
@@ -1517,6 +1549,9 @@
size_t global_work_offset[1];

global_work_offset[0] = work->blk.nonce;
+	if (opt_scrypt_jane)
+	applog(LOG_DEBUG, "Nonce: %x, Global work size: %x, local work size: %x", work->blk.nonce, (unsigned)globalThreads[0], (unsigned)localThreads[0]);
+
status = clEnqueueNDRangeKernel(clState->commandQueue, *kernel, 1, global_work_offset,
globalThreads, localThreads, 0, NULL, NULL);
} else
@@ -1534,6 +1569,12 @@
return -1;
}

+	if (opt_scrypt_jane) {
+	o = thrdata->res;
+	target = *(uint32_t *)(work->target + 28);
+	applog(LOG_DEBUG, "Nonce: %x, Output buffer: %x %x %x %x %x %x %x %x Target: %x", work->blk.nonce, o[0], o[1], o[2], o[3], o[4], o[5], o[6], o[7], target);
+	}
+
/* The amount of work scanned can fluctuate when intensity changes
* and since we do this one cycle behind, we increment the work more
* than enough to prevent repeating work */
diff -urN cgminer-3.7.2.orig/findnonce.c cgminer-3.7.2/findnonce.c
--- cgminer-3.7.2.orig/findnonce.c	2013-08-16 15:59:32.000000000 +1000
+++ cgminer-3.7.2/findnonce.c	2014-01-02 17:30:53.901239502 +1100
@@ -18,6 +18,7 @@

#include "findnonce.h"
#include "scrypt.h"
+#include "scrypt-jane.h"

const uint32_t SHA256_K[64] = {
0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5,
diff -urN cgminer-3.7.2.orig/Makefile.am cgminer-3.7.2/Makefile.am
--- cgminer-3.7.2.orig/Makefile.am	2013-11-06 00:24:07.000000000 +1100
+++ cgminer-3.7.2/Makefile.am	2014-01-02 17:30:53.905239579 +1100
@@ -54,7 +54,7 @@
cgminer_SOURCES += *.cl

if HAS_SCRYPT
-cgminer_SOURCES += scrypt.c scrypt.h
+cgminer_SOURCES += scrypt.c scrypt.h scrypt-jane.c scrypt-jane.h
endif

endif
diff -urN cgminer-3.7.2.orig/Makefile.in cgminer-3.7.2/Makefile.in
--- cgminer-3.7.2.orig/Makefile.in	2013-11-06 00:29:25.000000000 +1100
+++ cgminer-3.7.2/Makefile.in	2014-01-02 17:30:53.917239815 +1100
@@ -58,7 +58,7 @@
@HAS_OPENCL_TRUE@am__append_1 = driver-opencl.h driver-opencl.c ocl.c \
@HAS_OPENCL_TRUE@	ocl.h findnonce.c findnonce.h adl.c adl.h \
@HAS_OPENCL_TRUE@	adl_functions.h *.cl
-@HAS_OPENCL_TRUE@@HAS_SCRYPT_TRUE@am__append_2 = scrypt.c scrypt.h
+@HAS_OPENCL_TRUE@@HAS_SCRYPT_TRUE@am__append_2 = scrypt.c scrypt.h scrypt-jane.c scrypt-jane.h
@NEED_FPGAUTILS_TRUE@am__append_3 = fpgautils.c fpgautils.h
@WANT_USBUTILS_TRUE@am__append_4 = usbutils.c usbutils.h

@@ -110,18 +110,20 @@
bench_block.h util.c util.h uthash.h logging.h sha2.c sha2.h \
api.c logging.c driver-opencl.h driver-opencl.c ocl.c ocl.h \
findnonce.c findnonce.h adl.c adl.h adl_functions.h *.cl \
-	scrypt.c scrypt.h fpgautils.c fpgautils.h usbutils.c \
-	usbutils.h driver-avalon.c driver-avalon.h \
-	driver-knc-spi-fpga.c driver-bflsc.c driver-bflsc.h \
-	driver-bitforce.c driver-hashfast.c driver-hashfast.h \
-	hf_protocol.h driver-bitfury.c driver-bitfury.h \
-	driver-icarus.c driver-klondike.c driver-modminer.c
+	scrypt.c scrypt.h scrypt-jane.c scrypt-jane.h fpgautils.c \
+	fpgautils.h usbutils.c usbutils.h driver-avalon.c \
+	driver-avalon.h driver-knc-spi-fpga.c driver-bflsc.c \
+	driver-bflsc.h driver-bitforce.c driver-hashfast.c \
+	driver-hashfast.h hf_protocol.h driver-bitfury.c \
+	driver-bitfury.h driver-icarus.c driver-klondike.c \
+	driver-modminer.c
@HAS_OPENCL_TRUE@am__objects_1 = cgminer-driver-opencl.$(OBJEXT) \
@HAS_OPENCL_TRUE@	cgminer-ocl.$(OBJEXT) \
@HAS_OPENCL_TRUE@	cgminer-findnonce.$(OBJEXT) \
@HAS_OPENCL_TRUE@	cgminer-adl.$(OBJEXT)
@HAS_OPENCL_TRUE@@HAS_SCRYPT_TRUE@am__objects_2 = \
-@HAS_OPENCL_TRUE@@HAS_SCRYPT_TRUE@	cgminer-scrypt.$(OBJEXT)
+@HAS_OPENCL_TRUE@@HAS_SCRYPT_TRUE@	cgminer-scrypt.$(OBJEXT) \
+@HAS_OPENCL_TRUE@@HAS_SCRYPT_TRUE@	cgminer-scrypt-jane.$(OBJEXT)
@NEED_FPGAUTILS_TRUE@am__objects_3 = cgminer-fpgautils.$(OBJEXT)
@WANT_USBUTILS_TRUE@am__objects_4 = cgminer-usbutils.$(OBJEXT)
@HAS_AVALON_TRUE@am__objects_5 = cgminer-driver-avalon.$(OBJEXT)
@@ -735,6 +737,7 @@
@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/cgminer-fpgautils.Po@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/cgminer-logging.Po@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/cgminer-ocl.Po@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/cgminer-scrypt-jane.Po@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/cgminer-scrypt.Po@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/cgminer-sha2.Po@am__quote@
@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/cgminer-usbutils.Po@am__quote@
@@ -904,6 +907,20 @@
@AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
@am__fastdepCC_FALSE@	$(AM_V_CC@am__nodep@)$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(cgminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o cgminer-scrypt.obj `if test -f 'scrypt.c'; then $(CYGPATH_W) 'scrypt.c'; else $(CYGPATH_W) '$(srcdir)/scrypt.c'; fi`

+cgminer-scrypt-jane.o: scrypt-jane.c
+@am__fastdepCC_TRUE@	$(AM_V_CC)$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(cgminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT cgminer-scrypt-jane.o -MD -MP -MF $(DEPDIR)/cgminer-scrypt-jane.Tpo -c -o cgminer-scrypt-jane.o `test -f 'scrypt-jane.c' || echo '$(srcdir)/'`scrypt-jane.c
+@am__fastdepCC_TRUE@	$(AM_V_at)$(am__mv) $(DEPDIR)/cgminer-scrypt-jane.Tpo $(DEPDIR)/cgminer-scrypt-jane.Po
+@AMDEP_TRUE@@am__fastdepCC_FALSE@	$(AM_V_CC)source='scrypt-jane.c' object='cgminer-scrypt-jane.o' libtool=no @AMDEPBACKSLASH@
+@AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
+@am__fastdepCC_FALSE@	$(AM_V_CC@am__nodep@)$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(cgminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o cgminer-scrypt-jane.o `test -f 'scrypt-jane.c' || echo '$(srcdir)/'`scrypt-jane.c
+
+cgminer-scrypt-jane.obj: scrypt-jane.c
+@am__fastdepCC_TRUE@	$(AM_V_CC)$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(cgminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT cgminer-scrypt-jane.obj -MD -MP -MF $(DEPDIR)/cgminer-scrypt-jane.Tpo -c -o cgminer-scrypt-jane.obj `if test -f 'scrypt-jane.c'; then $(CYGPATH_W) 'scrypt-jane.c'; else $(CYGPATH_W) '$(srcdir)/scrypt-jane.c'; fi`
+@am__fastdepCC_TRUE@	$(AM_V_at)$(am__mv) $(DEPDIR)/cgminer-scrypt-jane.Tpo $(DEPDIR)/cgminer-scrypt-jane.Po
+@AMDEP_TRUE@@am__fastdepCC_FALSE@	$(AM_V_CC)source='scrypt-jane.c' object='cgminer-scrypt-jane.obj' libtool=no @AMDEPBACKSLASH@
+@AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
+@am__fastdepCC_FALSE@	$(AM_V_CC@am__nodep@)$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(cgminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o cgminer-scrypt-jane.obj `if test -f 'scrypt-jane.c'; then $(CYGPATH_W) 'scrypt-jane.c'; else $(CYGPATH_W) '$(srcdir)/scrypt-jane.c'; fi`
+
cgminer-fpgautils.o: fpgautils.c
@am__fastdepCC_TRUE@	$(AM_V_CC)$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(cgminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT cgminer-fpgautils.o -MD -MP -MF $(DEPDIR)/cgminer-fpgautils.Tpo -c -o cgminer-fpgautils.o `test -f 'fpgautils.c' || echo '$(srcdir)/'`fpgautils.c
@am__fastdepCC_TRUE@	$(AM_V_at)$(am__mv) $(DEPDIR)/cgminer-fpgautils.Tpo $(DEPDIR)/cgminer-fpgautils.Po
diff -urN cgminer-3.7.2.orig/miner.h cgminer-3.7.2/miner.h
--- cgminer-3.7.2.orig/miner.h	2013-11-06 00:24:07.000000000 +1100
+++ cgminer-3.7.2/miner.h	2014-01-02 17:30:53.925239975 +1100
@@ -405,6 +405,7 @@
KL_DIAKGCN,
KL_DIABLO,
KL_SCRYPT,
+	KL_SCRYPT_JANE,
};

enum dev_reason {
@@ -1089,6 +1090,11 @@
extern int opt_scantime;
extern int opt_expiry;

+/* scrypt-jane */
+extern unsigned int sj_minNf;
+extern unsigned int sj_maxNf;
+extern unsigned int sj_startTime;
+
#ifdef USE_USBUTILS
extern pthread_mutex_t cgusb_lock;
extern pthread_mutex_t cgusbres_lock;
@@ -1171,6 +1177,7 @@
extern int gpu_threads;
#ifdef USE_SCRYPT
extern bool opt_scrypt;
+extern bool opt_scrypt_jane;
#else
#define opt_scrypt (0)
#endif
@@ -1612,4 +1619,6 @@
extern struct api_data *api_add_diff(struct api_data *root, char *name, double *data, bool copy_data);
extern struct api_data *api_add_percent(struct api_data *root, char *name, double *data, bool copy_data);

+extern unsigned char sj_GetNfactor(int nTimestamp);
+
#endif /* __MINER_H__ */
diff -urN cgminer-3.7.2.orig/ocl.c cgminer-3.7.2/ocl.c
--- cgminer-3.7.2.orig/ocl.c	2013-10-15 19:35:16.000000000 +1100
+++ cgminer-3.7.2/ocl.c	2014-01-02 17:30:53.933240131 +1100
@@ -392,7 +392,13 @@
if (cgpu->kernel == KL_NONE) {
if (opt_scrypt) {
applog(LOG_INFO, "Selecting scrypt kernel");
-	clState->chosen_kernel = KL_SCRYPT;
+	if (opt_scrypt_jane) {
+	applog(LOG_INFO, "Selected scrypt-jane kernel");
+	clState->chosen_kernel = KL_SCRYPT_JANE;
+	} else {
+	applog(LOG_INFO, "Selected standard scrypt kernel");
+	clState->chosen_kernel = KL_SCRYPT;
+	}
} else if (!strstr(name, "Tahiti") &&
/* Detect all 2.6 SDKs not with Tahiti and use diablo kernel */
(strstr(vbuff, "844.4") || // Linux 64 bit ATI 2.6 SDK
@@ -455,6 +461,12 @@
/* Scrypt only supports vector 1 */
cgpu->vwidth = 1;
break;
+	case KL_SCRYPT_JANE:
+	strcpy(filename, SCRYPT_JANE_KERNNAME".cl");
+	strcpy(binaryfilename, SCRYPT_JANE_KERNNAME);
+	/* Scrypt only supports vector 1 */
+	cgpu->vwidth = 1;
+	break;
case KL_NONE: /* Shouldn't happen */
case KL_DIABLO:
strcpy(filename, DIABLO_KERNNAME".cl");
diff -urN cgminer-3.7.2.orig/scrypt-jane.c cgminer-3.7.2/scrypt-jane.c
--- cgminer-3.7.2.orig/scrypt-jane.c	1970-01-01 10:00:00.000000000 +1000
+++ cgminer-3.7.2/scrypt-jane.c	2014-01-02 17:31:18.173719101 +1100
@@ -0,0 +1,703 @@
+/*-
+ * Copyright 2009 Colin Percival, 2011 ArtForz
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ * 1. Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ * 2. Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE AUTHOR AND CONTRIBUTORS ``AS IS'' AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED. IN NO EVENT SHALL THE AUTHOR OR CONTRIBUTORS BE LIABLE
+ * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+ * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS
+ * OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION)
+ * HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
+ * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY
+ * OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF
+ * SUCH DAMAGE.
+ *
+ * This file was originally written by Colin Percival as part of the Tarsnap
+ * online backup system.
+ */
+
+#include "config.h"
+#include "miner.h"
+
+#include
+#include
+#include
+
+
+// yacoin: increasing Nfactor gradually
+//const unsigned char minNfactor = 4;
+//const unsigned char maxNfactor = 30;
+
+unsigned char sj_GetNfactor(int nTimestamp) {
+ int l = 0;
+
+ if (nTimestamp <= sj_startTime)
+ return sj_minNf;
+
+ int s = nTimestamp - sj_startTime;
+ while ((s >> 1) > 3) {
+ l += 1;
+ s >>= 1;
+ }
+
+ s &= 3;
+
+ int n = (l * 170 + s * 25 - 2320) / 100;
+
+ if (n < 0) n = 0;
+
+ if (n > 255)
+ printf("sj_GetNfactor(%d) - something wrong(n == %d)\n", nTimestamp, n);
+
+ unsigned char N = (unsigned char)n;
+
+ if(Nsj_maxNf) return sj_maxNf;
+ return N;
+}
+
+/*
+ * Encode a length len/4 vector of (uint32_t) into a length len vector of
+ * (unsigned char) in big-endian form. Assumes len is a multiple of 4.
+ */
+inline void
+sj_be32enc_vect(uint32_t *dst, const uint32_t *src, uint32_t len)
+{
+	uint32_t i;
+
+	for (i = 0; i < len; i++)
+	dst[i] = htobe32(src[i]);
+}
+
+static void sj_scrypt(const uint8_t *password, size_t password_len, const uint8_t *salt, size_t salt_len, uint8_t Nfactor, uint8_t rfactor, uint8_t pfactor, uint8_t *out, size_t bytes);
+
+void sj_scrypt_regenhash(struct work *work)
+{
+	uint32_t data[20];
+	uint32_t *nonce = (uint32_t *)(work->data + 76);
+	uint32_t *ohash = (uint32_t *)(work->hash);
+
+	sj_be32enc_vect(data, (const uint32_t *)work->data, 19);
+	data[19] = htobe32(*nonce);
+
+	applog(LOG_DEBUG, "timestamp %d", data[17]);
+
+ int nfactor = sj_GetNfactor(data[17]);
+
+ applog(LOG_DEBUG, "Dat0: %x %x %x %x %x %x %x %x %x %x %x %x %x %x %x %x %x %x %x %x", data[0], data[1], data[2], data[3], data[4], data[5], data[6], data[7], data[8], data[9],
+ data[10], data[11], data[12], data[13], data[14], data[15], data[16], data[17], data[18], data[19]);
+
+ sj_scrypt((unsigned char *)data, 80,
+ (unsigned char *)data, 80,
+ nfactor, 0, 0, (unsigned char *)ohash, 32);
+
+//	flip32(ohash, ohash); // Not needed for scrypt-chacha - mikaelh
+ uint32_t *o = ohash;
+	applog(LOG_DEBUG, "Nonce: %x, Output buffe0: %x %x %x %x %x %x %x %x", *nonce, o[0], o[1], o[2], o[3], o[4], o[5], o[6], o[7]);
+}
+
+static const uint32_t sj_diff1targ = 0x0000ffff;
+
+typedef uint32_t sj_scrypt_mix_word_t;
+
+typedef void (*sj_scrypt_fatal_errorfn)(const char *msg);
+void sj_scrypt_set_fatal_error(sj_scrypt_fatal_errorfn fn);
+
+#define SJ_SCRYPT_BLOCK_BYTES 64
+#define SJ_SCRYPT_BLOCK_WORDS (SJ_SCRYPT_BLOCK_BYTES / sizeof(sj_scrypt_mix_word_t))
+
+#define SJ_ROTL32(a,b) (((a) << (b)) | ((a) >> (32 - b)))
+#define SJ_ROTL64(a,b) (((a) << (b)) | ((a) >> (64 - b)))
+#define SJ_U8TO64_LE(p) \
+	(((uint64_t)SJ_U8TO32_LE(p)) | ((uint64_t)SJ_U8TO32_LE((p) + 4) << 32))
+#define SJ_U8TO32_LE(p) \
+	(((uint32_t)((p)[0]) ) | ((uint32_t)((p)[1]) << 8) | \
+	((uint32_t)((p)[2]) << 16) | ((uint32_t)((p)[3]) << 24))
+#define SJ_U32TO8_BE(p, v) \
+	(p)[0] = (uint8_t)((v) >> 24); (p)[1] = (uint8_t)((v) >> 16); \
+	(p)[2] = (uint8_t)((v) >> 8); (p)[3] = (uint8_t)((v) );
+#define SJ_U64TO8_LE(p, v) \
+	SJ_U32TO8_LE((p), (uint32_t)((v) )); \
+	SJ_U32TO8_LE((p) + 4, (uint32_t)((v) >> 32));
+#define SJ_U32TO8_LE(p, v) \
+	(p)[0] = (uint8_t)((v) ); (p)[1] = (uint8_t)((v) >> 8); \
+	(p)[2] = (uint8_t)((v) >> 16); (p)[3] = (uint8_t)((v) >> 24);
+
+#define sj_scrypt_maxN 30 /* (1 << (30 + 1)) = ~2 billion */
+#if (SJ_SCRYPT_BLOCK_BYTES == 64)
+#define sj_scrypt_r_32kb 8 /* (1 << 8) = 256 * 2 blocks in a chunk * 64 bytes = Max of 32kb in a chunk */
+#elif (SJ_SCRYPT_BLOCK_BYTES == 128)
+#define sj_scrypt_r_32kb 7 /* (1 << 7) = 128 * 2 blocks in a chunk * 128 bytes = Max of 32kb in a chunk */
+#elif (SJ_SCRYPT_BLOCK_BYTES == 256)
+#define sj_scrypt_r_32kb 6 /* (1 << 6) = 64 * 2 blocks in a chunk * 256 bytes = Max of 32kb in a chunk */
+#elif (SJ_SCRYPT_BLOCK_BYTES == 512)
+#define sj_scrypt_r_32kb 5 /* (1 << 5) = 32 * 2 blocks in a chunk * 512 bytes = Max of 32kb in a chunk */
+#endif
+#define sj_scrypt_maxr sj_scrypt_r_32kb /* 32kb */
+#define sj_scrypt_maxp 25 /* (1 << 25) = ~33 million */
+
+#include
+#include
+
+static void
+sj_scrypt_fatal_error_default(const char *msg) {
+	fprintf(stderr, "%s\n", msg);
+	exit(1);
+}
+
+static sj_scrypt_fatal_errorfn sj_scrypt_fatal_error = sj_scrypt_fatal_error_default;
+
+void
+sj_scrypt_set_fatal_error_default(sj_scrypt_fatal_errorfn fn) {
+	sj_scrypt_fatal_error = fn;
+}
+
+typedef struct sj_scrypt_aligned_alloc_t {
+	uint8_t *mem, *ptr;
+} sj_scrypt_aligned_alloc;
+
+#if defined(SCRYPT_TEST_SPEED)
+static uint8_t *sj_mem_base = (uint8_t *)0;
+static size_t sj_mem_bump = 0;
+
+/* allocations are assumed to be multiples of 64 bytes and total allocations not to exceed ~1.01gb */
+static sj_scrypt_aligned_alloc
+sj_scrypt_alloc(uint64_t size) {
+	sj_scrypt_aligned_alloc aa;
+	if (!sj_mem_base) {
+	sj_mem_base = (uint8_t *)malloc((1024 * 1024 * 1024) + (1024 * 1024) + (SJ_SCRYPT_BLOCK_BYTES - 1));
+	if (!sj_mem_base)
+	sj_scrypt_fatal_error("scrypt-jane: out of memory");
+	sj_mem_base = (uint8_t *)(((size_t)sj_mem_base + (SJ_SCRYPT_BLOCK_BYTES - 1)) & ~(SJ_SCRYPT_BLOCK_BYTES - 1));
+	}
+	aa.mem = sj_mem_base + sj_mem_bump;
+	aa.ptr = aa.mem;
+	sj_mem_bump += (size_t)size;
+	return aa;
+}
+
+static void
+sj_scrypt_free(sj_scrypt_aligned_alloc *aa) {
+	sj_mem_bump = 0;
+}
+#else
+static sj_scrypt_aligned_alloc
+sj_scrypt_alloc(uint64_t size) {
+	static const size_t max_alloc = (size_t)-1;
+	sj_scrypt_aligned_alloc aa;
+	size += (SJ_SCRYPT_BLOCK_BYTES - 1);
+	if (size > max_alloc)
+	sj_scrypt_fatal_error("scrypt-jane: not enough address space on this CPU to allocate required memory");
+	aa.mem = (uint8_t *)malloc((size_t)size);
+	aa.ptr = (uint8_t *)(((size_t)aa.mem + (SJ_SCRYPT_BLOCK_BYTES - 1)) & ~(SJ_SCRYPT_BLOCK_BYTES - 1));
+	if (!aa.mem)
+	sj_scrypt_fatal_error("scrypt-jane: out of memory");
+	return aa;
+}
+
+static void
+sj_scrypt_free(sj_scrypt_aligned_alloc *aa) {
+	free(aa->mem);
+}
+#endif
+
+
+
+static void
+sj_chacha_core(uint32_t state[16]) {
+	size_t rounds = 8;
+	uint32_t x0,x1,x2,x3,x4,x5,x6,x7,x8,x9,x10,x11,x12,x13,x14,x15,t;
+
+	x0 = state[0];
+	x1 = state[1];
+	x2 = state[2];
+	x3 = state[3];
+	x4 = state[4];
+	x5 = state[5];
+	x6 = state[6];
+	x7 = state[7];
+	x8 = state[8];
+	x9 = state[9];
+	x10 = state[10];
+	x11 = state[11];
+	x12 = state[12];
+	x13 = state[13];
+	x14 = state[14];
+	x15 = state[15];
+
+	#define quarter(a,b,c,d) \
+	a += b; t = d^a; d = SJ_ROTL32(t,16); \
+	c += d; t = b^c; b = SJ_ROTL32(t,12); \
+	a += b; t = d^a; d = SJ_ROTL32(t, 8); \
+	c += d; t = b^c; b = SJ_ROTL32(t, 7);
+
+	for (; rounds; rounds -= 2) {
+	quarter( x0, x4, x8,x12)
+	quarter( x1, x5, x9,x13)
+	quarter( x2, x6,x10,x14)
+	quarter( x3, x7,x11,x15)
+	quarter( x0, x5,x10,x15)
+	quarter( x1, x6,x11,x12)
+	quarter( x2, x7, x8,x13)
+	quarter( x3, x4, x9,x14)
+	}
+
+	state[0] += x0;
+	state[1] += x1;
+	state[2] += x2;
+	state[3] += x3;
+	state[4] += x4;
+	state[5] += x5;
+	state[6] += x6;
+	state[7] += x7;
+	state[8] += x8;
+	state[9] += x9;
+	state[10] += x10;
+	state[11] += x11;
+	state[12] += x12;
+	state[13] += x13;
+	state[14] += x14;
+	state[15] += x15;
+
+	#undef quarter
+}
+
+
+/* returns a pointer to item i, where item is len sj_scrypt_mix_word_t's long */
+static sj_scrypt_mix_word_t *
+sj_scrypt_item(sj_scrypt_mix_word_t *base, sj_scrypt_mix_word_t i, sj_scrypt_mix_word_t len) {
+	return base + (i * len);
+}
+
+
+static sj_scrypt_mix_word_t *
+sj_scrypt_block(sj_scrypt_mix_word_t *base, sj_scrypt_mix_word_t i) {
+	return base + (i * SJ_SCRYPT_BLOCK_WORDS);
+}
+
+#define SJ_MM16 __attribute__((aligned(16)))
+
+static void
+sj_scrypt_ChunkMix(sj_scrypt_mix_word_t *Bout/*[chunkWords]*/, sj_scrypt_mix_word_t *Bin/*[chunkWords]*/, sj_scrypt_mix_word_t *Bxor/*[chunkWords]*/, uint32_t r) {
+	sj_scrypt_mix_word_t SJ_MM16 X[SJ_SCRYPT_BLOCK_WORDS], *block;
+	uint32_t i, j, blocksPerChunk = r * 2, half = 0;
+
+	/* 1: X = B_{2r - 1} */
+	block = sj_scrypt_block(Bin, blocksPerChunk - 1);
+	for (i = 0; i < SJ_SCRYPT_BLOCK_WORDS; i++)
+	X[i] = block[i];
+
+	if (Bxor) {
+	block = sj_scrypt_block(Bxor, blocksPerChunk - 1);
+	for (i = 0; i < SJ_SCRYPT_BLOCK_WORDS; i++)
+	X[i] ^= block[i];
+	}
+
+	/* 2: for i = 0 to 2r - 1 do */
+	for (i = 0; i < blocksPerChunk; i++, half ^= r) {
+	/* 3: X = H(X ^ B_i) */
+	block = sj_scrypt_block(Bin, i);
+	for (j = 0; j < SJ_SCRYPT_BLOCK_WORDS; j++)
+	X[j] ^= block[j];
+
+	if (Bxor) {
+	block = sj_scrypt_block(Bxor, i);
+	for (j = 0; j < SJ_SCRYPT_BLOCK_WORDS; j++)
+	X[j] ^= block[j];
+	}
+	/* SCRYPT_MIX_FN */ sj_chacha_core(X);
+
+	/* 4: Y_i = X */
+	/* 6: B'[0..r-1] = Y_even */
+	/* 6: B'[r..2r-1] = Y_odd */
+	block = sj_scrypt_block(Bout, (i / 2) + half);
+	for (j = 0; j < SJ_SCRYPT_BLOCK_WORDS; j++)
+	block[j] = X[j];
+	}
+}
+
+#define SJ_U32_SWAP(v) { \
+	(v) = (((v) << 8) & 0xFF00FF00 ) | (((v) >> 8) & 0xFF00FF ); \
+ (v) = ((v) << 16) | ((v) >> 16); \
+}
+
+#define SJ_SCRYPT_WORD_ENDIAN_SWAP SJ_U32_SWAP
+
+/* romix pre/post endian conversion function */
+static void
+sj_scrypt_romix_convert_endian(sj_scrypt_mix_word_t *blocks, size_t nblocks) {
+#if !defined(CPU_LE)
+	static const union { uint8_t b[2]; uint16_t w; } endian_test = {{1,0}};
+	size_t i;
+	if (endian_test.w == 0x100) {
+	nblocks *= SJ_SCRYPT_BLOCK_WORDS;
+	for (i = 0; i < nblocks; i++) {
+	SJ_SCRYPT_WORD_ENDIAN_SWAP(blocks[i]);
+	}
+	}
+#endif
+}
+
+static void
+sj_scrypt_ROMix(sj_scrypt_mix_word_t *X/*[chunkWords]*/, sj_scrypt_mix_word_t *Y/*[chunkWords]*/, sj_scrypt_mix_word_t *V/*[N * chunkWords]*/, uint32_t N, uint32_t r) {
+	uint32_t i, j, chunkWords = SJ_SCRYPT_BLOCK_WORDS * r * 2;
+	sj_scrypt_mix_word_t *block = V;
+
+	sj_scrypt_romix_convert_endian(X, r * 2);
+
+	/* 1: X = B */
+	/* implicit */
+
+	/* 2: for i = 0 to N - 1 do */
+	memcpy(block, X, chunkWords * sizeof(sj_scrypt_mix_word_t));
+	for (i = 0; i < N - 1; i++, block += chunkWords) {
+	/* 3: V_i = X */
+	/* 4: X = H(X) */
+	sj_scrypt_ChunkMix(block + chunkWords, block, NULL, r);
+	}
+	sj_scrypt_ChunkMix(X, block, NULL, r);
+
+	/* 6: for i = 0 to N - 1 do */
+	for (i = 0; i < N; i += 2) {
+	/* 7: j = Integerify(X) % N */
+	j = X[chunkWords - SJ_SCRYPT_BLOCK_WORDS] & (N - 1);
+
+	/* 8: X = H(Y ^ V_j) */
+	sj_scrypt_ChunkMix(Y, X, sj_scrypt_item(V, j, chunkWords), r);
+
+	/* 7: j = Integerify(Y) % N */
+	j = Y[chunkWords - SJ_SCRYPT_BLOCK_WORDS] & (N - 1);
+
+	/* 8: X = H(Y ^ V_j) */
+	sj_scrypt_ChunkMix(X, Y, sj_scrypt_item(V, j, chunkWords), r);
+	}
+
+	/* 10: B' = X */
+	/* implicit */
+
+	sj_scrypt_romix_convert_endian(X, r * 2);
+}
+
+#define SJ_SCRYPT_HASH "Keccak-512"
+#define SJ_SCRYPT_HASH_DIGEST_SIZE 64
+#define SJ_SCRYPT_KECCAK_F 1600
+#define SJ_SCRYPT_KECCAK_C (SJ_SCRYPT_HASH_DIGEST_SIZE * 8 * 2) /* 256=512, 512=1024 */
+#define SJ_SCRYPT_KECCAK_R (SJ_SCRYPT_KECCAK_F - SJ_SCRYPT_KECCAK_C) /* 256=1088, 512=576 */
+#define SJ_SCRYPT_HASH_BLOCK_SIZE (SJ_SCRYPT_KECCAK_R / 8)
+
+typedef uint8_t sj_scrypt_hash_digest[SJ_SCRYPT_HASH_DIGEST_SIZE];
+
+typedef struct sj_scrypt_hash_state_t {
+	uint64_t state[SJ_SCRYPT_KECCAK_F / 64];
+	uint32_t leftover;
+	uint8_t buffer[SJ_SCRYPT_HASH_BLOCK_SIZE];
+} sj_scrypt_hash_state;
+
+typedef struct sj_scrypt_hmac_state_t {
+	sj_scrypt_hash_state inner, outer;
+} sj_scrypt_hmac_state;
+
+static const uint64_t sj_keccak_round_constants[24] = {
+	0x0000000000000001ull, 0x0000000000008082ull,
+	0x800000000000808aull, 0x8000000080008000ull,
+	0x000000000000808bull, 0x0000000080000001ull,
+	0x8000000080008081ull, 0x8000000000008009ull,
+	0x000000000000008aull, 0x0000000000000088ull,
+	0x0000000080008009ull, 0x000000008000000aull,
+	0x000000008000808bull, 0x800000000000008bull,
+	0x8000000000008089ull, 0x8000000000008003ull,
+	0x8000000000008002ull, 0x8000000000000080ull,
+	0x000000000000800aull, 0x800000008000000aull,
+	0x8000000080008081ull, 0x8000000000008080ull,
+	0x0000000080000001ull, 0x8000000080008008ull
+};
+
+static void
+sj_keccak_block(sj_scrypt_hash_state *S, const uint8_t *in) {
+	size_t i;
+	uint64_t *s = S->state, t[5], u[5], v, w;
+
+	/* absorb input */
+	for (i = 0; i < SJ_SCRYPT_HASH_BLOCK_SIZE / 8; i++, in += 8)
+	s[i] ^= SJ_U8TO64_LE(in);
+
+	for (i = 0; i < 24; i++) {
+	/* theta: c = a[0,i] ^ a[1,i] ^ .. a[4,i] */
+	t[0] = s[0] ^ s[5] ^ s[10] ^ s[15] ^ s[20];
+	t[1] = s[1] ^ s[6] ^ s[11] ^ s[16] ^ s[21];
+	t[2] = s[2] ^ s[7] ^ s[12] ^ s[17] ^ s[22];
+	t[3] = s[3] ^ s[8] ^ s[13] ^ s[18] ^ s[23];
+	t[4] = s[4] ^ s[9] ^ s[14] ^ s[19] ^ s[24];
+
+	/* theta: d[i] = c[i+4] ^ rotl(c[i+1],1) */
+	u[0] = t[4] ^ SJ_ROTL64(t[1], 1);
+	u[1] = t[0] ^ SJ_ROTL64(t[2], 1);
+	u[2] = t[1] ^ SJ_ROTL64(t[3], 1);
+	u[3] = t[2] ^ SJ_ROTL64(t[4], 1);
+	u[4] = t[3] ^ SJ_ROTL64(t[0], 1);
+
+	/* theta: a[0,i], a[1,i], .. a[4,i] ^= d[i] */
+	s[0] ^= u[0]; s[5] ^= u[0]; s[10] ^= u[0]; s[15] ^= u[0]; s[20] ^= u[0];
+	s[1] ^= u[1]; s[6] ^= u[1]; s[11] ^= u[1]; s[16] ^= u[1]; s[21] ^= u[1];
+	s[2] ^= u[2]; s[7] ^= u[2]; s[12] ^= u[2]; s[17] ^= u[2]; s[22] ^= u[2];
+	s[3] ^= u[3]; s[8] ^= u[3]; s[13] ^= u[3]; s[18] ^= u[3]; s[23] ^= u[3];
+	s[4] ^= u[4]; s[9] ^= u[4]; s[14] ^= u[4]; s[19] ^= u[4]; s[24] ^= u[4];
+
+	/* rho pi: b[..] = rotl(a[..], ..) */
+	v = s[ 1];
+	s[ 1] = SJ_ROTL64(s[ 6], 44);
+	s[ 6] = SJ_ROTL64(s[ 9], 20);
+	s[ 9] = SJ_ROTL64(s[22], 61);
+	s[22] = SJ_ROTL64(s[14], 39);
+	s[14] = SJ_ROTL64(s[20], 18);
+	s[20] = SJ_ROTL64(s[ 2], 62);
+	s[ 2] = SJ_ROTL64(s[12], 43);
+	s[12] = SJ_ROTL64(s[13], 25);
+	s[13] = SJ_ROTL64(s[19], 8);
+	s[19] = SJ_ROTL64(s[23], 56);
+	s[23] = SJ_ROTL64(s[15], 41);
+	s[15] = SJ_ROTL64(s[ 4], 27);
+	s[ 4] = SJ_ROTL64(s[24], 14);
+	s[24] = SJ_ROTL64(s[21], 2);
+	s[21] = SJ_ROTL64(s[ 8], 55);
+	s[ 8] = SJ_ROTL64(s[16], 45);
+	s[16] = SJ_ROTL64(s[ 5], 36);
+	s[ 5] = SJ_ROTL64(s[ 3], 28);
+	s[ 3] = SJ_ROTL64(s[18], 21);
+	s[18] = SJ_ROTL64(s[17], 15);
+	s[17] = SJ_ROTL64(s[11], 10);
+	s[11] = SJ_ROTL64(s[ 7], 6);
+	s[ 7] = SJ_ROTL64(s[10], 3);
+	s[10] = SJ_ROTL64( v, 1);
+
+	/* chi: a[i,j] ^= ~b[i,j+1] & b[i,j+2] */
+	v = s[ 0]; w = s[ 1]; s[ 0] ^= (~w) & s[ 2]; s[ 1] ^= (~s[ 2]) & s[ 3]; s[ 2] ^= (~s[ 3]) & s[ 4]; s[ 3] ^= (~s[ 4]) & v; s[ 4] ^= (~v) & w;
+	v = s[ 5]; w = s[ 6]; s[ 5] ^= (~w) & s[ 7]; s[ 6] ^= (~s[ 7]) & s[ 8]; s[ 7] ^= (~s[ 8]) & s[ 9]; s[ 8] ^= (~s[ 9]) & v; s[ 9] ^= (~v) & w;
+	v = s[10]; w = s[11]; s[10] ^= (~w) & s[12]; s[11] ^= (~s[12]) & s[13]; s[12] ^= (~s[13]) & s[14]; s[13] ^= (~s[14]) & v; s[14] ^= (~v) & w;
+	v = s[15]; w = s[16]; s[15] ^= (~w) & s[17]; s[16] ^= (~s[17]) & s[18]; s[17] ^= (~s[18]) & s[19]; s[18] ^= (~s[19]) & v; s[19] ^= (~v) & w;
+	v = s[20]; w = s[21]; s[20] ^= (~w) & s[22]; s[21] ^= (~s[22]) & s[23]; s[22] ^= (~s[23]) & s[24]; s[23] ^= (~s[24]) & v; s[24] ^= (~v) & w;
+
+	/* iota: a[0,0] ^= round constant */
+	s[0] ^= sj_keccak_round_constants[i];
+	}
+}
+
+static void
+sj_scrypt_hash_init(sj_scrypt_hash_state *S) {
+	memset(S, 0, sizeof(*S));
+}
+
+static void
+sj_scrypt_hash_update(sj_scrypt_hash_state *S, const uint8_t *in, size_t inlen) {
+	size_t want;
+
+	/* handle the previous data */
+	if (S->leftover) {
+	want = (SJ_SCRYPT_HASH_BLOCK_SIZE - S->leftover);
+	want = (want < inlen) ? want : inlen;
+	memcpy(S->buffer + S->leftover, in, want);
+	S->leftover += (uint32_t)want;
+	if (S->leftover < SJ_SCRYPT_HASH_BLOCK_SIZE)
+	return;
+	in += want;
+	inlen -= want;
+	sj_keccak_block(S, S->buffer);
+	}
+
+	/* handle the current data */
+	while (inlen >= SJ_SCRYPT_HASH_BLOCK_SIZE) {
+	sj_keccak_block(S, in);
+	in += SJ_SCRYPT_HASH_BLOCK_SIZE;
+	inlen -= SJ_SCRYPT_HASH_BLOCK_SIZE;
+	}
+
+	/* handle leftover data */
+	S->leftover = (uint32_t)inlen;
+	if (S->leftover)
+	memcpy(S->buffer, in, S->leftover);
+}
+
+static void
+sj_scrypt_hash_finish(sj_scrypt_hash_state *S, uint8_t *hash) {
+	size_t i;
+
+	S->buffer[S->leftover] = 0x01;
+	memset(S->buffer + (S->leftover + 1), 0, SJ_SCRYPT_HASH_BLOCK_SIZE - (S->leftover + 1));
+	S->buffer[SJ_SCRYPT_HASH_BLOCK_SIZE - 1] |= 0x80;
+	sj_keccak_block(S, S->buffer);
+
+	for (i = 0; i < SJ_SCRYPT_HASH_DIGEST_SIZE; i += 8) {
+	SJ_U64TO8_LE(&hash[i], S->state[i / 8]);
+	}
+}
+
+static void
+sj_scrypt_hash(sj_scrypt_hash_digest hash, const uint8_t *m, size_t mlen) {
+	sj_scrypt_hash_state st;
+	sj_scrypt_hash_init(&st);
+	sj_scrypt_hash_update(&st, m, mlen);
+	sj_scrypt_hash_finish(&st, hash);
+}
+
+/* hmac */
+static void
+sj_scrypt_hmac_init(sj_scrypt_hmac_state *st, const uint8_t *key, size_t keylen) {
+	uint8_t pad[SJ_SCRYPT_HASH_BLOCK_SIZE] = {0};
+	size_t i;
+
+	sj_scrypt_hash_init(&st->inner);
+	sj_scrypt_hash_init(&st->outer);
+
+	if (keylen <= SJ_SCRYPT_HASH_BLOCK_SIZE) {
+	/* use the key directly if it's <= blocksize bytes */
+	memcpy(pad, key, keylen);
+	} else {
+	/* if it's > blocksize bytes, hash it */
+	sj_scrypt_hash(pad, key, keylen);
+	}
+
+	/* inner = (key ^ 0x36) */
+	/* h(inner || ...) */
+	for (i = 0; i < SJ_SCRYPT_HASH_BLOCK_SIZE; i++)
+	pad[i] ^= 0x36;
+	sj_scrypt_hash_update(&st->inner, pad, SJ_SCRYPT_HASH_BLOCK_SIZE);
+
+	/* outer = (key ^ 0x5c) */
+	/* h(outer || ...) */
+	for (i = 0; i < SJ_SCRYPT_HASH_BLOCK_SIZE; i++)
+	pad[i] ^= (0x5c ^ 0x36);
+	sj_scrypt_hash_update(&st->outer, pad, SJ_SCRYPT_HASH_BLOCK_SIZE);
+}
+
+static void
+sj_scrypt_hmac_update(sj_scrypt_hmac_state *st, const uint8_t *m, size_t mlen) {
+	/* h(inner || m...) */
+	sj_scrypt_hash_update(&st->inner, m, mlen);
+}
+
+static void
+sj_scrypt_hmac_finish(sj_scrypt_hmac_state *st, sj_scrypt_hash_digest mac) {
+	/* h(inner || m) */
+	sj_scrypt_hash_digest innerhash;
+	sj_scrypt_hash_finish(&st->inner, innerhash);
+
+	/* h(outer || h(inner || m)) */
+	sj_scrypt_hash_update(&st->outer, innerhash, sizeof(innerhash));
+	sj_scrypt_hash_finish(&st->outer, mac);
+}
+
+static void
+sj_scrypt_pbkdf2(const uint8_t *password, size_t password_len, const uint8_t *salt, size_t salt_len, uint8_t *out, size_t bytes) {
+	sj_scrypt_hmac_state hmac_pw, hmac_pw_salt, work;
+	sj_scrypt_hash_digest ti;
+	uint8_t be[4];
+	uint32_t i, blocks;
+
+	/* bytes must be <= (0xffffffff - (SCRYPT_HASH_DIGEST_SIZE - 1)), which they will always be under scrypt */
+
+	/* hmac(password, ...) */
+	sj_scrypt_hmac_init(&hmac_pw, password, password_len);
+
+	/* hmac(password, salt...) */
+	hmac_pw_salt = hmac_pw;
+	sj_scrypt_hmac_update(&hmac_pw_salt, salt, salt_len);
+
+	blocks = ((uint32_t)bytes + (SJ_SCRYPT_HASH_DIGEST_SIZE - 1)) / SJ_SCRYPT_HASH_DIGEST_SIZE;
+	for (i = 1; i <= blocks; i++) {
+	/* U1 = hmac(password, salt || be(i)) */
+	SJ_U32TO8_BE(be, i);
+	work = hmac_pw_salt;
+	sj_scrypt_hmac_update(&work, be, 4);
+	sj_scrypt_hmac_finish(&work, ti);
+
+	/* T[i] = U1 ^ U2 ^ U3... */
+
+	memcpy(out, ti, (bytes > SJ_SCRYPT_HASH_DIGEST_SIZE) ? SJ_SCRYPT_HASH_DIGEST_SIZE : bytes);
+	out += SJ_SCRYPT_HASH_DIGEST_SIZE;
+	bytes -= SJ_SCRYPT_HASH_DIGEST_SIZE;
+	}
+}
+
+
+static void
+sj_scrypt(const uint8_t *password, size_t password_len, const uint8_t *salt, size_t salt_len, uint8_t Nfactor, uint8_t rfactor, uint8_t pfactor, uint8_t *out, size_t bytes) {
+	sj_scrypt_aligned_alloc YX, V;
+	uint8_t *X, *Y;
+	uint32_t N, r, p, chunk_bytes;
+
+	if (Nfactor > sj_scrypt_maxN)
+	sj_scrypt_fatal_error("scrypt-jane: N out of range");
+	if (rfactor > sj_scrypt_maxr)
+	sj_scrypt_fatal_error("scrypt-jane: r out of range");
+	if (pfactor > sj_scrypt_maxp)
+	sj_scrypt_fatal_error("scrypt-jane: p out of range");
+
+	N = (1 << (Nfactor + 1));
+	r = (1 << rfactor);
+	p = (1 << pfactor);
+
+	chunk_bytes = SJ_SCRYPT_BLOCK_BYTES * r * 2;
+	V = sj_scrypt_alloc((uint64_t)N * chunk_bytes);
+	YX = sj_scrypt_alloc((p + 1) * chunk_bytes);
+
+	/* 1: X = PBKDF2(password, salt) */
+	Y = YX.ptr;
+	X = Y + chunk_bytes;
+ sj_scrypt_pbkdf2(password, password_len, salt, salt_len, X, chunk_bytes);
+
+	/* 2: X = ROMix(X) */
+	sj_scrypt_ROMix((sj_scrypt_mix_word_t *)X, (sj_scrypt_mix_word_t *)Y, (sj_scrypt_mix_word_t *)V.ptr, N, 1);
+
+	/* 3: Out = PBKDF2(password, X) */
+	sj_scrypt_pbkdf2(password, password_len, X, chunk_bytes, out, bytes);
+
+	sj_scrypt_free(&V);
+	sj_scrypt_free(&YX);
+}
+
+/* Used externally as confirmation of correct OCL code */
+int sj_scrypt_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t nonce)
+{
+	uint32_t tmp_hash7, Htarg = le32toh(((const uint32_t *)ptarget)[7]);
+	uint32_t data[20], ohash[8];
+
+	sj_be32enc_vect(data, (const uint32_t *)pdata, 19);
+	data[19] = nonce;
+
+ int nfactor = sj_GetNfactor(data[17]);
+
+ applog(LOG_DEBUG, "Dat0: %x %x %x %x %x %x %x %x %x %x %x %x %x %x %x %x %x %x %x %x", data[0], data[1], data[2], data[3], data[4], data[5], data[6], data[7], data[8], data[9],
+ data[10], data[11], data[12], data[13], data[14], data[15], data[16], data[17], data[18], data[19]);
+
+ sj_scrypt((unsigned char *)data, 80,
+ (unsigned char *)data, 80,
+ nfactor, 0, 0, (unsigned char *)ohash, 32);
+
+ uint32_t *o = ohash;
+	applog(LOG_DEBUG, "Nonce: %x, Output buffe0: %x %x %x %x %x %x %x %x", nonce, o[0], o[1], o[2], o[3], o[4], o[5], o[6], o[7]);
+
+	tmp_hash7 = ohash[7];
+
+	applog(LOG_DEBUG, "Nonce %x harget %08lx diff1 %08lx hash %08lx",
+ nonce,
+	(long unsigned int)Htarg,
+	(long unsigned int)sj_diff1targ,
+	(long unsigned int)tmp_hash7);
+	if (tmp_hash7 > sj_diff1targ)
+	return -1;
+	if (tmp_hash7 > Htarg)
+	return 0;
+	return 1;
+}
+
diff -urN cgminer-3.7.2.orig/scrypt-jane.cl cgminer-3.7.2/scrypt-jane.cl
--- cgminer-3.7.2.orig/scrypt-jane.cl	1970-01-01 10:00:00.000000000 +1000
+++ cgminer-3.7.2/scrypt-jane.cl	2014-01-02 17:30:53.945240369 +1100
@@ -0,0 +1,705 @@
+/*
+	scrypt-jane by Andrew M, https://github.com/floodyberry/scrypt-jane
+
+	OpenCL version by hanzac
+
+	Support for LOOKUP_GAP and CONCURRENT_THREADS by mikaelh
+	Nfactor compensation by mikaelh
+	Keccak rewrite by mikaelh
+
+	Public Domain or MIT License, whichever is easier
+*/
+
+#define SCRYPT_HASH "Keccak-512"
+#define SCRYPT_HASH_DIGEST_SIZE 64
+#define SCRYPT_KECCAK_F 1600
+#define SCRYPT_HASH_BLOCK_SIZE 72
+#define SCRYPT_BLOCK_BYTES 128
+#define ROTL64(x, y) as_uint2(rotate(as_ulong(x), y))
+#define ROTL32(x, y) rotate(x, y)
+
+typedef struct scrypt_hash_state_t {
+	uint4 state4[(SCRYPT_KECCAK_F + 127) / 128]; // 8 bytes of extra
+	uint4 buffer4[(SCRYPT_HASH_BLOCK_SIZE + 15) / 16]; // 8 bytes of extra
+	//uint leftover;
+} scrypt_hash_state;
+
+typedef struct scrypt_hmac_state_t {
+	scrypt_hash_state inner;
+	scrypt_hash_state outer;
+} scrypt_hmac_state;
+
+
+__constant ulong keccak_round_constants[24] = {
+	0x0000000000000001UL, 0x0000000000008082UL,
+	0x800000000000808aUL, 0x8000000080008000UL,
+	0x000000000000808bUL, 0x0000000080000001UL,
+	0x8000000080008081UL, 0x8000000000008009UL,
+	0x000000000000008aUL, 0x0000000000000088UL,
+	0x0000000080008009UL, 0x000000008000000aUL,
+	0x000000008000808bUL, 0x800000000000008bUL,
+	0x8000000000008089UL, 0x8000000000008003UL,
+	0x8000000000008002UL, 0x8000000000000080UL,
+	0x000000000000800aUL, 0x800000008000000aUL,
+	0x8000000080008081UL, 0x8000000000008080UL,
+	0x0000000080000001UL, 0x8000000080008008UL
+};
+
+
+static void
+keccak_block_core(scrypt_hash_state *S) {
+	uint2 t[5];
+	uint2 u[5];
+	uint2 v;
+	uint2 w;
+	uint4 *s4 = S->state4;
+	uint i;
+
+	for (i = 0; i < 24; i++) {
+	/* theta: c = a[0,i] ^ a[1,i] ^ .. a[4,i] */
+	t[0] = s4[0].xy ^ s4[2].zw ^ s4[5].xy ^ s4[7].zw ^ s4[10].xy;
+	t[1] = s4[0].zw ^ s4[3].xy ^ s4[5].zw ^ s4[8].xy ^ s4[10].zw;
+	t[2] = s4[1].xy ^ s4[3].zw ^ s4[6].xy ^ s4[8].zw ^ s4[11].xy;
+	t[3] = s4[1].zw ^ s4[4].xy ^ s4[6].zw ^ s4[9].xy ^ s4[11].zw;
+	t[4] = s4[2].xy ^ s4[4].zw ^ s4[7].xy ^ s4[9].zw ^ s4[12].xy;
+
+	/* theta: d[i] = c[i+4] ^ rotl(c[i+1],1) */
+	u[0] = t[4] ^ ROTL64(t[1], 1UL);
+	u[1] = t[0] ^ ROTL64(t[2], 1UL);
+	u[2] = t[1] ^ ROTL64(t[3], 1UL);
+	u[3] = t[2] ^ ROTL64(t[4], 1UL);
+	u[4] = t[3] ^ ROTL64(t[0], 1UL);
+
+	/* theta: a[0,i], a[1,i], .. a[4,i] ^= d[i] */
+	s4[0].xy ^= u[0]; s4[2].zw ^= u[0]; s4[5].xy ^= u[0]; s4[7].zw ^= u[0]; s4[10].xy ^= u[0];
+	s4[0].zw ^= u[1]; s4[3].xy ^= u[1]; s4[5].zw ^= u[1]; s4[8].xy ^= u[1]; s4[10].zw ^= u[1];
+	s4[1].xy ^= u[2]; s4[3].zw ^= u[2]; s4[6].xy ^= u[2]; s4[8].zw ^= u[2]; s4[11].xy ^= u[2];
+	s4[1].zw ^= u[3]; s4[4].xy ^= u[3]; s4[6].zw ^= u[3]; s4[9].xy ^= u[3]; s4[11].zw ^= u[3];
+	s4[2].xy ^= u[4]; s4[4].zw ^= u[4]; s4[7].xy ^= u[4]; s4[9].zw ^= u[4]; s4[12].xy ^= u[4];
+
+	/* rho pi: b[..] = rotl(a[..], ..) */
+	v = s4[0].zw;
+	s4[ 0].zw = ROTL64(s4[ 3].xy, 44UL);
+	s4[ 3].xy = ROTL64(s4[ 4].zw, 20UL);
+	s4[ 4].zw = ROTL64(s4[11].xy, 61UL);
+	s4[11].xy = ROTL64(s4[ 7].xy, 39UL);
+	s4[ 7].xy = ROTL64(s4[10].xy, 18UL);
+	s4[10].xy = ROTL64(s4[ 1].xy, 62UL);
+	s4[ 1].xy = ROTL64(s4[ 6].xy, 43UL);
+	s4[ 6].xy = ROTL64(s4[ 6].zw, 25UL);
+	s4[ 6].zw = ROTL64(s4[ 9].zw, 8UL);
+	s4[ 9].zw = ROTL64(s4[11].zw, 56UL);
+	s4[11].zw = ROTL64(s4[ 7].zw, 41UL);
+	s4[ 7].zw = ROTL64(s4[ 2].xy, 27UL);
+	s4[ 2].xy = ROTL64(s4[12].xy, 14UL);
+	s4[12].xy = ROTL64(s4[10].zw, 2UL);
+	s4[10].zw = ROTL64(s4[ 4].xy, 55UL);
+	s4[ 4].xy = ROTL64(s4[ 8].xy, 45UL);
+	s4[ 8].xy = ROTL64(s4[ 2].zw, 36UL);
+	s4[ 2].zw = ROTL64(s4[ 1].zw, 28UL);
+	s4[ 1].zw = ROTL64(s4[ 9].xy, 21UL);
+	s4[ 9].xy = ROTL64(s4[ 8].zw, 15UL);
+	s4[ 8].zw = ROTL64(s4[ 5].zw, 10UL);
+	s4[ 5].zw = ROTL64(s4[ 3].zw, 6UL);
+	s4[ 3].zw = ROTL64(s4[ 5].xy, 3UL);
+	s4[ 5].xy = ROTL64( v, 1UL);
+
+	/* chi: a[i,j] ^= ~b[i,j+1] & b[i,j+2] */
+	v = s4[ 0].xy; w = s4[ 0].zw; s4[ 0].xy ^= (~w) & s4[ 1].xy; s4[ 0].zw ^= (~s4[ 1].xy) & s4[ 1].zw; s4[ 1].xy ^= (~s4[ 1].zw) & s4[ 2].xy; s4[ 1].zw ^= (~s4[ 2].xy) & v; s4[ 2].xy ^= (~v) & w;
+	v = s4[ 2].zw; w = s4[ 3].xy; s4[ 2].zw ^= (~w) & s4[ 3].zw; s4[ 3].xy ^= (~s4[ 3].zw) & s4[ 4].xy; s4[ 3].zw ^= (~s4[ 4].xy) & s4[ 4].zw; s4[ 4].xy ^= (~s4[ 4].zw) & v; s4[ 4].zw ^= (~v) & w;
+	v = s4[ 5].xy; w = s4[ 5].zw; s4[ 5].xy ^= (~w) & s4[ 6].xy; s4[ 5].zw ^= (~s4[ 6].xy) & s4[ 6].zw; s4[ 6].xy ^= (~s4[ 6].zw) & s4[ 7].xy; s4[ 6].zw ^= (~s4[ 7].xy) & v; s4[ 7].xy ^= (~v) & w;
+	v = s4[ 7].zw; w = s4[ 8].xy; s4[ 7].zw ^= (~w) & s4[ 8].zw; s4[ 8].xy ^= (~s4[ 8].zw) & s4[ 9].xy; s4[ 8].zw ^= (~s4[ 9].xy) & s4[ 9].zw; s4[ 9].xy ^= (~s4[ 9].zw) & v; s4[ 9].zw ^= (~v) & w;
+	v = s4[10].xy; w = s4[10].zw; s4[10].xy ^= (~w) & s4[11].xy; s4[10].zw ^= (~s4[11].xy) & s4[11].zw; s4[11].xy ^= (~s4[11].zw) & s4[12].xy; s4[11].zw ^= (~s4[12].xy) & v; s4[12].xy ^= (~v) & w;
+
+	/* iota: a[0,0] ^= round constant */
+	s4[0].xy ^= as_uint2(keccak_round_constants[i]);
+	}
+}
+
+__constant uint4 ZERO = (uint4)(0);
+__constant uint2 ZERO_UINT2 = (uint2)(0);
+
+static void
+keccak_block(scrypt_hash_state *S, const uint4 *in4) {
+	uint4 *s4 = S->state4;
+	uint i;
+
+	/* absorb input */
+	#pragma unroll
+	for (i = 0; i < 4; i++) {
+	s4[i] ^= in4[i];
+	}
+	s4[4].xy ^= in4[4].xy;
+
+	keccak_block_core(S);
+}
+
+static void
+keccak_block_zero(scrypt_hash_state *S, const uint4 *in4) {
+	uint4 *s4 = S->state4;
+	uint i;
+
+	/* absorb input */
+	#pragma unroll
+	for (i = 0; i < 4; i++) {
+	s4[i] = in4[i];
+	}
+	s4[4].xyzw = (uint4)(in4[4].xy, 0, 0);
+
+	#pragma unroll
+	for (i = 5; i < 12; i++) {
+	s4[i] = ZERO;
+	}
+	s4[12].xy = ZERO_UINT2;
+
+	keccak_block_core(S);
+}
+
+static void
+scrypt_hash_update_72(scrypt_hash_state *S, const uint4 *in4) {
+	/* handle the current data */
+	keccak_block_zero(S, in4);
+}
+
+static void
+scrypt_hash_update_80(scrypt_hash_state *S, const uint4 *in4) {
+	const uchar *in = (const uchar *)in4;
+	uint i;
+
+	/* handle the current data */
+	keccak_block(S, in4);
+	in += SCRYPT_HASH_BLOCK_SIZE;
+
+	/* handle leftover data */
+	//S->leftover = 2;
+
+	{
+	const uint2 *int2 = (const uint2 *) in;
+
+	S->buffer4[0].xy = int2[0].xy;
+	}
+}
+
+static void
+scrypt_hash_update_128(scrypt_hash_state *S, const uint4 *in4) {
+	const uchar *in = (const uchar *)in4;
+	uint i;
+
+	/* handle the current data */
+	keccak_block(S, in4);
+	in += SCRYPT_HASH_BLOCK_SIZE;
+
+	/* handle leftover data */
+	//S->leftover = 14;
+
+	{
+	const uint2 *int2 = (const uint2 *) in;
+
+	#pragma unroll
+	for (i = 0; i < 3; i++) {
+	S->buffer4[i] = (uint4)(int2[2 * i].xy, int2[2 * i + 1].xy);
+	}
+	S->buffer4[3].xy = int2[6].xy;
+	}
+}
+
+static void
+scrypt_hash_update_4_after_80(scrypt_hash_state *S, uint in) {
+	// assume that leftover = 2
+	/* handle the previous data */
+	S->buffer4[0].zw = (uint2)(in, 0x01);
+	//S->leftover += 1;
+}
+
+static void
+scrypt_hash_update_4_after_128(scrypt_hash_state *S, uint in) {
+	// leftover = 14
+	/* handle the previous data */
+	S->buffer4[3].zw = (uint2)(in, 0x01);
+	//S->leftover += 1;
+}
+
+static void
+scrypt_hash_update_64(scrypt_hash_state *S, const uint4 *in4) {
+	uint i;
+
+	/* handle leftover data */
+	//S->leftover = 16;
+
+	#pragma unroll
+	for (i = 0; i < 4; i++) {
+	S->buffer4[i] = in4[i];
+	}
+}
+
+static void
+scrypt_hash_finish_80_after_64(scrypt_hash_state *S, uint4 *hash4) {
+	// assume that leftover = 16
+	S->buffer4[4].xy = (uint2)(0x01, 0x80000000);
+
+	keccak_block(S, S->buffer4);
+
+	#pragma unroll
+	for (uint i = 0; i < 4; i++) {
+	hash4[i] = S->state4[i];
+	}
+}
+
+static void
+scrypt_hash_finish_80_after_80_4(scrypt_hash_state *S, uint4 *hash4) {
+	uint i;
+	// assume that leftover = 3
+	//S->buffer4[0].w = 0x01; // done already in scrypt_hash_update_4_after_80
+	#pragma unroll
+	for (i = 1; i < 4; i++) {
+	S->buffer4[i] = ZERO;
+	}
+	S->buffer4[4].xy = (uint2)(0, 0x80000000);
+
+	keccak_block(S, S->buffer4);
+
+	#pragma unroll
+	for (uint i = 0; i < 4; i++) {
+	hash4[i] = S->state4[i];
+	}
+}
+
+static void
+scrypt_hash_finish_80_after_128_4(scrypt_hash_state *S, uint4 *hash4) {
+	// leftover = 15
+	//S->buffer4[3].w = 0x01; // done already in scrypt_hash_update_4_after_128
+	S->buffer4[4].xy = (uint2)(0, 0x80000000);
+
+	keccak_block(S, S->buffer4);
+
+	#pragma unroll
+	for (uint i = 0; i < 4; i++) {
+	hash4[i] = S->state4[i];
+	}
+}
+
+static void
+scrypt_hash_80(uint4 *hash4, const uint4 *m) {
+	const uchar *in = (const uchar *)m;
+	scrypt_hash_state st;
+	uint i;
+
+	/* handle the current data */
+	keccak_block_zero(&st, m);
+	in += SCRYPT_HASH_BLOCK_SIZE;
+
+	{
+	const uint2 *in2 = (const uint2 *) in;
+	st.buffer4[0].xyzw = (uint4)(in2[0].xy, 0x01, 0);
+	}
+
+	#pragma unroll
+	for (i = 1; i < 4; i++) {
+	st.buffer4[i] = ZERO;
+	}
+	st.buffer4[4].xyzw = (uint4)(0, 0x80000000, 0, 0);
+
+	keccak_block(&st, st.buffer4);
+
+	#pragma unroll
+	for (i = 0; i < 4; i++) {
+	hash4[i] = st.state4[i];
+	}
+}
+
+/* hmac */
+__constant uint4 KEY_0X36 = (uint4)(0x36363636);
+__constant uint2 KEY_0X36_2 = (uint2)(0x36363636);
+__constant uint4 KEY_0X36_XOR_0X5C = (uint4)(0x6A6A6A6A);
+__constant uint2 KEY_0X36_XOR_0X5C_2 = (uint2)(0x6A6A6A6A);
+
+static void
+scrypt_hmac_init(scrypt_hmac_state *st, const uint4 *key) {
+	uint4 pad4[SCRYPT_HASH_BLOCK_SIZE/16 + 1];
+	uint i;
+
+	/* if it's > blocksize bytes, hash it */
+	scrypt_hash_80(pad4, key);
+	pad4[4].xy = ZERO_UINT2;
+
+	/* inner = (key ^ 0x36) */
+	/* h(inner || ...) */
+	#pragma unroll
+	for (i = 0; i < 4; i++) {
+	pad4[i] ^= KEY_0X36;
+	}
+	pad4[4].xy ^= KEY_0X36_2;
+	scrypt_hash_update_72(&st->inner, pad4);
+
+	/* outer = (key ^ 0x5c) */
+	/* h(outer || ...) */
+	#pragma unroll
+	for (i = 0; i < 4; i++) {
+	pad4[i] ^= KEY_0X36_XOR_0X5C;
+	}
+	pad4[4].xy ^= KEY_0X36_XOR_0X5C_2;
+	scrypt_hash_update_72(&st->outer, pad4);
+}
+
+static void
+scrypt_hmac_update_80(scrypt_hmac_state *st, const uint4 *m) {
+	/* h(inner || m...) */
+	scrypt_hash_update_80(&st->inner, m);
+}
+
+static void
+scrypt_hmac_update_128(scrypt_hmac_state *st, const uint4 *m) {
+	/* h(inner || m...) */
+	scrypt_hash_update_128(&st->inner, m);
+}
+
+static void
+scrypt_hmac_update_4_after_80(scrypt_hmac_state *st, uint m) {
+	/* h(inner || m...) */
+	scrypt_hash_update_4_after_80(&st->inner, m);
+}
+
+static void
+scrypt_hmac_update_4_after_128(scrypt_hmac_state *st, uint m) {
+	/* h(inner || m...) */
+	scrypt_hash_update_4_after_128(&st->inner, m);
+}
+
+static void
+scrypt_hmac_finish_128B(scrypt_hmac_state *st, uint4 *mac) {
+	/* h(inner || m) */
+	uint4 innerhash[4];
+	scrypt_hash_finish_80_after_80_4(&st->inner, innerhash);
+
+	/* h(outer || h(inner || m)) */
+	scrypt_hash_update_64(&st->outer, innerhash);
+	scrypt_hash_finish_80_after_64(&st->outer, mac);
+}
+
+static void
+scrypt_hmac_finish_32B(scrypt_hmac_state *st, uint4 *mac) {
+	/* h(inner || m) */
+	uint4 innerhash[4];
+	scrypt_hash_finish_80_after_128_4(&st->inner, innerhash);
+
+	/* h(outer || h(inner || m)) */
+	scrypt_hash_update_64(&st->outer, innerhash);
+	scrypt_hash_finish_80_after_64(&st->outer, mac);
+}
+
+static void
+scrypt_copy_hmac_state_128B(scrypt_hmac_state *dest, const scrypt_hmac_state *src) {
+	uint i;
+
+	for (i = 0; i < 12; i++) {
+	dest->inner.state4[i] = src->inner.state4[i];
+	}
+	dest->inner.state4[12].xy = src->inner.state4[12].xy;
+
+	dest->inner.buffer4[0].xy = src->inner.buffer4[0].xy;
+
+	for (i = 0; i < 12; i++) {
+	dest->outer.state4[i] = src->outer.state4[i];
+	}
+	dest->outer.state4[12].xy = src->outer.state4[12].xy;
+}
+
+__constant uint be1 = 0x01000000;
+__constant uint be2 = 0x02000000;
+
+static void
+scrypt_pbkdf2_128B(const uint4 *password, const uint4 *salt, uint4 *out4) {
+	scrypt_hmac_state hmac_pw, work;
+	uint4 ti4[4];
+	uint i;
+
+	/* bytes must be <= (0xffffffff - (SCRYPT_HASH_DIGEST_SIZE - 1)), which they will always be under scrypt */
+
+	/* hmac(password, ...) */
+	scrypt_hmac_init(&hmac_pw, password);
+
+	/* hmac(password, salt...) */
+	scrypt_hmac_update_80(&hmac_pw, salt);
+
+	/* U1 = hmac(password, salt || be(i)) */
+	/* U32TO8_BE(be, i); */
+	//work = hmac_pw;
+	scrypt_copy_hmac_state_128B(&work, &hmac_pw);
+	scrypt_hmac_update_4_after_80(&work, be1);
+	scrypt_hmac_finish_128B(&work, ti4);
+
+	#pragma unroll
+	for (i = 0; i < 4; i++) {
+	out4[i] = ti4[i];
+	}
+
+	/* U1 = hmac(password, salt || be(i)) */
+	/* U32TO8_BE(be, i); */
+	// work = hmac_pw;
+	scrypt_hmac_update_4_after_80(&hmac_pw, be2);
+	scrypt_hmac_finish_128B(&hmac_pw, ti4);
+
+	#pragma unroll
+	for (i = 0; i < 4; i++) {
+	out4[i + 4] = ti4[i];
+	}
+}
+
+static void
+scrypt_pbkdf2_32B(const uint4 *password, const uint4 *salt, uint4 *out4) {
+	scrypt_hmac_state hmac_pw;
+	uint4 ti4[4];
+
+	/* bytes must be <= (0xffffffff - (SCRYPT_HASH_DIGEST_SIZE - 1)), which they will always be under scrypt */
+
+	/* hmac(password, ...) */
+	scrypt_hmac_init(&hmac_pw, password);
+
+	/* hmac(password, salt...) */
+	scrypt_hmac_update_128(&hmac_pw, salt);
+
+	/* U1 = hmac(password, salt || be(i)) */
+	/* U32TO8_BE(be, i); */
+	scrypt_hmac_update_4_after_128(&hmac_pw, be1);
+	scrypt_hmac_finish_32B(&hmac_pw, ti4);
+
+	#pragma unroll
+	for (uint i = 0; i < 2; i++) {
+	out4[i] = ti4[i];
+	}
+}
+
+__constant uint4 MASK_2 = (uint4) (1, 2, 3, 0);
+__constant uint4 MASK_3 = (uint4) (2, 3, 0, 1);
+__constant uint4 MASK_4 = (uint4) (3, 0, 1, 2);
+__constant uint4 ROTATE_16 = (uint4) (16, 16, 16, 16);
+__constant uint4 ROTATE_12 = (uint4) (12, 12, 12, 12);
+__constant uint4 ROTATE_8 = (uint4) (8, 8, 8, 8);
+__constant uint4 ROTATE_7 = (uint4) (7, 7, 7, 7);
+
+static void
+chacha_core(uint4 state[4]) {
+	uint4 x[4];
+	uint4 t;
+	uint rounds;
+
+	x[0] = state[0];
+	x[1] = state[1];
+	x[2] = state[2];
+	x[3] = state[3];
+
+	#pragma unroll
+	for (rounds = 0; rounds < 4; rounds ++) {
+	x[0] += x[1]; t = x[3] ^ x[0]; x[3] = ROTL32(t, ROTATE_16);
+	x[2] += x[3]; t = x[1] ^ x[2]; x[1] = ROTL32(t, ROTATE_12);
+	x[0] += x[1]; t = x[3] ^ x[0]; x[3] = ROTL32(t, ROTATE_8);
+	x[2] += x[3]; t = x[1] ^ x[2]; x[1] = ROTL32(t, ROTATE_7);
+
+	// x[1] = shuffle(x[1], MASK_2);
+	// x[2] = shuffle(x[2], MASK_3);
+	// x[3] = shuffle(x[3], MASK_4);
+
+	x[0] += x[1].yzwx; t = x[3].wxyz ^ x[0]; x[3].wxyz = ROTL32(t, ROTATE_16);
+	x[2].zwxy += x[3].wxyz; t = x[1].yzwx ^ x[2].zwxy; x[1].yzwx = ROTL32(t, ROTATE_12);
+	x[0] += x[1].yzwx; t = x[3].wxyz ^ x[0]; x[3].wxyz = ROTL32(t, ROTATE_8);
+	x[2].zwxy += x[3].wxyz; t = x[1].yzwx ^ x[2].zwxy; x[1].yzwx = ROTL32(t, ROTATE_7);
+
+	// x[1] = shuffle(x[1], MASK_4);
+	// x[2] = shuffle(x[2], MASK_3);
+	// x[3] = shuffle(x[3], MASK_2);
+	}
+
+	state[0] += x[0];
+	state[1] += x[1];
+	state[2] += x[2];
+	state[3] += x[3];
+}
+
+static void
+scrypt_ChunkMix_inplace_Bxor_local(uint4 *restrict B/*[chunkWords]*/, uint4 *restrict Bxor/*[chunkWords]*/) {
+	/* 1: X = B_{2r - 1} */
+
+	/* 2: for i = 0 to 2r - 1 do */
+	/* 3: X = H(X ^ B_i) */
+	B[0] ^= B[4] ^ Bxor[4] ^ Bxor[0];
+	B[1] ^= B[5] ^ Bxor[5] ^ Bxor[1];
+	B[2] ^= B[6] ^ Bxor[6] ^ Bxor[2];
+	B[3] ^= B[7] ^ Bxor[7] ^ Bxor[3];
+
+	/* SCRYPT_MIX_FN */ chacha_core(B);
+
+	/* 4: Y_i = X */
+	/* 6: B'[0..r-1] = Y_even */
+	/* 6: B'[r..2r-1] = Y_odd */
+
+
+	/* 3: X = H(X ^ B_i) */
+	B[4] ^= B[0] ^ Bxor[4];
+	B[5] ^= B[1] ^ Bxor[5];
+	B[6] ^= B[2] ^ Bxor[6];
+	B[7] ^= B[3] ^ Bxor[7];
+
+	/* SCRYPT_MIX_FN */ chacha_core(B + 4);
+
+	/* 4: Y_i = X */
+	/* 6: B'[0..r-1] = Y_even */
+	/* 6: B'[r..2r-1] = Y_odd */
+}
+
+static void
+scrypt_ChunkMix_inplace_local(uint4 *restrict B/*[chunkWords]*/) {
+	/* 1: X = B_{2r - 1} */
+
+	/* 2: for i = 0 to 2r - 1 do */
+	/* 3: X = H(X ^ B_i) */
+	B[0] ^= B[4];
+	B[1] ^= B[5];
+	B[2] ^= B[6];
+	B[3] ^= B[7];
+
+	/* SCRYPT_MIX_FN */ chacha_core(B);
+
+	/* 4: Y_i = X */
+	/* 6: B'[0..r-1] = Y_even */
+	/* 6: B'[r..2r-1] = Y_odd */
+
+
+	/* 3: X = H(X ^ B_i) */
+	B[4] ^= B[0];
+	B[5] ^= B[1];
+	B[6] ^= B[2];
+	B[7] ^= B[3];
+
+	/* SCRYPT_MIX_FN */ chacha_core(B + 4);
+
+	/* 4: Y_i = X */
+	/* 6: B'[0..r-1] = Y_even */
+	/* 6: B'[r..2r-1] = Y_odd */
+}
+
+#define Coord(x,y,z) x+y*(x ## SIZE)+z*(y ## SIZE)*(x ## SIZE)
+#define CO Coord(z,x,y)
+
+static void
+scrypt_ROMix(uint4 *restrict X/*[chunkWords]*/, __global uint4 *restrict lookup/*[N * chunkWords]*/, const uint N, const uint gid, const uint Nfactor) {
+	const uint effective_concurrency = (CONCURRENT_THREADS << 9) >> Nfactor;
+	const uint zSIZE = 8;
+	const uint ySIZE = (N/LOOKUP_GAP+(N%LOOKUP_GAP>0));
+	const uint xSIZE = effective_concurrency;
+	const uint x = gid % xSIZE;
+	uint i, j, y, z;
+	uint4 W[8];
+
+	/* 1: X = B */
+	/* implicit */
+
+	/* 2: for i = 0 to N - 1 do */
+	for (y = 0; y < N / LOOKUP_GAP; y++) {
+	/* 3: V_i = X */
+	#pragma unroll
+	for (z = 0; z < zSIZE; z++) {
+	lookup[CO] = X[z];
+	}
+
+	for (j = 0; j < LOOKUP_GAP; j++) {
+	/* 4: X = H(X) */
+	scrypt_ChunkMix_inplace_local(X);
+	}
+	}
+
+#if (LOOKUP_GAP != 1) && (LOOKUP_GAP != 2) && (LOOKUP_GAP != 4) && (LOOKUP_GAP != 8)
+	if (N % LOOKUP_GAP > 0) {
+	y = N / LOOKUP_GAP;
+
+	#pragma unroll
+	for (z = 0; z < zSIZE; z++) {
+	lookup[CO] = X[z];
+	}
+
+	for (j = 0; j < N % LOOKUP_GAP; j++) {
+	scrypt_ChunkMix_inplace_local(X);
+	}
+	}
+#endif
+
+	/* 6: for i = 0 to N - 1 do */
+	for (i = 0; i < N; i++) {
+	/* 7: j = Integerify(X) % N */
+	j = X[4].x & (N - 1);
+	y = j / LOOKUP_GAP;
+
+	#pragma unroll
+	for (z = 0; z < zSIZE; z++) {
+	W[z] = lookup[CO];
+	}
+
+
+#if (LOOKUP_GAP == 1)
+#elif (LOOKUP_GAP == 2)
+	if (j & 1) {
+	scrypt_ChunkMix_inplace_local(W);
+	}
+#else
+	uint c = j % LOOKUP_GAP;
+	for (uint k = 0; k < c; k++) {
+	scrypt_ChunkMix_inplace_local(W);
+	}
+#endif
+
+	/* 8: X = H(X ^ V_j) */
+	scrypt_ChunkMix_inplace_Bxor_local(X, W);
+	}
+
+	/* 10: B' = X */
+	/* implicit */
+}
+
+__constant uint ES[2] = { 0x00FF00FF, 0xFF00FF00 };
+#define FOUND (0xFF)
+#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
+#define EndianSwap(n) (rotate(n & Es2[0].x, 24U)|rotate(n & Es2[0].y, 8U))
+
+__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
+__kernel void search(__global const uint4 * restrict input,
+volatile __global uint * restrict output, __global uchar * restrict padcache,
+const uint4 midstate0, const uint4 midstate16, const uint target, const uint N)
+{
+	uint4 password[5];
+	uint4 X[8];
+	uint output_hash[8] __attribute__ ((aligned (16)));
+	const uint gid = get_global_id(0);
+	uint Nfactor = 0;
+	uint tmp = N >> 1;
+
+	/* Shortcut if Nfactor is at least 9 which it currently is */
+	if ((tmp & 512 - 1) == 0) {
+	Nfactor += 9;
+	tmp >>= 9;
+	}
+
+	/* Determine the Nfactor */
+	while ((tmp & 1) == 0) {
+	tmp >>= 1;
+	Nfactor++;
+	}
+
+	password[0] = input[0];
+	password[1] = input[1];
+	password[2] = input[2];
+	password[3] = input[3];
+	password[4] = input[4];
+	password[4].w = gid;
+
+	/* 1: X = PBKDF2(password, salt) */
+	scrypt_pbkdf2_128B(password, password, X);
+
+	/* 2: X = ROMix(X) */
+	scrypt_ROMix(X, (__global uint4 *)padcache, N, gid, Nfactor);
+
+	/* 3: Out = PBKDF2(password, X) */
+	scrypt_pbkdf2_32B(password, X, (uint4 *)output_hash);
+
+	bool result = (output_hash[7] <= target);
+	if (result)
+	SETFOUND(gid);
+}
diff -urN cgminer-3.7.2.orig/scrypt-jane.h cgminer-3.7.2/scrypt-jane.h
--- cgminer-3.7.2.orig/scrypt-jane.h	1970-01-01 10:00:00.000000000 +1000
+++ cgminer-3.7.2/scrypt-jane.h	2014-01-02 17:31:08.505527949 +1100
@@ -0,0 +1,26 @@
+#ifndef SCRYPT_JANE_H
+#define SCRYPT_JANE_H
+
+#include "miner.h"
+
+#ifdef USE_SCRYPT
+extern int sj_scrypt_test(unsigned char *pdata, const unsigned char *ptarget,
+	uint32_t nonce);
+extern void sj_scrypt_regenhash(struct work *work);
+
+extern inline void sj_be32enc_vect(uint32_t *dst, const uint32_t *src, uint32_t len);
+
+#else /* USE_SCRYPT */
+static inline int sj_scrypt_test(__maybe_unused unsigned char *pdata,
+	__maybe_unused const unsigned char *ptarget,
+	__maybe_unused uint32_t nonce)
+{
+	return 0;
+}
+
+static inline void sj_scrypt_regenhash(__maybe_unused struct work *work)
+{
+}
+#endif /* USE_SCRYPT */
+
+#endif /* SCRYPT_JANE_H */