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

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_JANEKERNNAME + / 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_be32encvect / +#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 = bswap32(((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@amappend_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@amappend_2 = scrypt.c scrypt.h +@HAS_OPENCL_TRUE@@HAS_SCRYPT_TRUE@amappend_2 = scrypt.c scrypt.h scrypt-jane.c scrypt-jane.h @NEED_FPGAUTILS_TRUE@am__append_3 = fpgautils.c fpgautils.h @WANT_USBUTILS_TRUE@amappend_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@amobjects_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@amobjects_3 = cgminer-fpgautils.$(OBJEXT) @WANT_USBUTILS_TRUE@amobjects_4 = cgminer-usbutils.$(OBJEXT) @HAS_AVALON_TRUE@amobjects_5 = cgminer-driver-avalon.$(OBJEXT) @@ -735,6 +737,7 @@ @AMDEP_TRUE@@aminclude@ @amquote@./$(DEPDIR)/cgminer-fpgautils.Po@amquote@ @AMDEP_TRUE@@aminclude@ @amquote@./$(DEPDIR)/cgminer-logging.Po@amquote@ @AMDEP_TRUE@@aminclude@ @amquote@./$(DEPDIR)/cgminer-ocl.Po@amquote@ +@AMDEP_TRUE@@aminclude@ @amquote@./$(DEPDIR)/cgminer-scrypt-jane.Po@am__quote@ @AMDEP_TRUE@@aminclude@ @amquote@./$(DEPDIR)/cgminer-scrypt.Po@amquote@ @AMDEP_TRUE@@aminclude@ @amquote@./$(DEPDIR)/cgminer-sha2.Po@amquote@ @AMDEP_TRUE@@aminclude@ @amquote@./$(DEPDIR)/cgminer-usbutils.Po@am__quote@ @@ -904,6 +907,20 @@ @AMDEP_TRUE@@amfastdepCC_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 +@amfastdepCC_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 +@amfastdepCC_TRUE@ $(AM_V_at)$(ammv) $(DEPDIR)/cgminer-scrypt-jane.Tpo $(DEPDIR)/cgminer-scrypt-jane.Po +@AMDEP_TRUE@@amfastdepCC_FALSE@ $(AM_V_CC)source='scrypt-jane.c' object='cgminer-scrypt-jane.o' libtool=no @AMDEPBACKSLASH@ +@AMDEP_TRUE@@amfastdepCC_FALSE@ DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@ +@amfastdepCC_FALSE@ $(AM_V_CC@amnodep@)$(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 +@amfastdepCC_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 +@amfastdepCC_TRUE@ $(AM_V_at)$(ammv) $(DEPDIR)/cgminer-scrypt-jane.Tpo $(DEPDIR)/cgminer-scrypt-jane.Po +@AMDEP_TRUE@@amfastdepCC_FALSE@ $(AM_V_CC)source='scrypt-jane.c' object='cgminer-scrypt-jane.obj' libtool=no @AMDEPBACKSLASH@ +@AMDEP_TRUE@@amfastdepCC_FALSE@ DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@ +@amfastdepCC_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 @amfastdepCC_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 KLNONE: /* Shouldn't happen / case KL_DIABLO: strcpy(filename, DIABLOKERNNAME".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 @@ +/-
*
* 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(N<sj_minNf) return sj_minNf;
if(N>sj_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_scryptmaxN 30 /* (1 << (30 + 1)) = ~2 billion / +#if (SJ_SCRYPT_BLOCK_BYTES == 64) +#define sj_scrypt_r32kb 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_r32kb 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_r32kb 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_r32kb 5 / (1 << 5) = 32 * 2 blocks in a chunk * 512 bytes = Max of 32kb in a chunk _/ +#endif +#define sj_scrypt_maxr sj_scrypt_r32kb / 32kb _/ +#define sj_scryptmaxp 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_membump = 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_wordt *Bin/[chunkWords]_/, sj_scrypt_mix_wordt *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_wordt *Y/[chunkWords]_/, sj_scrypt_mix_wordt *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_DIGESTSIZE * 8 * 2) /* 256=512, 512=1024 / +#define SJ_SCRYPT_KECCAK_R (SJ_SCRYPT_KECCAK_F - SJ_SCRYPT_KECCAKC) / 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] = Yodd / +} + +static void +scrypt_ChunkMix_inplacelocal(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] = Yodd / +} + +#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, uint32t 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 / USESCRYPT / + +#endif /_ SCRYPT_JANE_H */

jgarzik / cpuminer

add scypt-jane support to 3.7.2 #56

include "driver-opencl.h"

include "bench_block.h"

include "scrypt.h"

ifdef USE_USBUTILS

include "usbutils.h"

endif

if defined(HAVE_OPENCL) || defined(USE_USBUTILS)

endif

ifdef USE_SCRYPT

endif

endif

ifdef USE_AVALON

ifdef USE_SCRYPT

undef SCRYPT_KERNNAME

undef SIG_ATOMIC_T_SUFFIX

include "adl.h"

include "util.h"

ifdef HAVE_CURSES

ifdef USE_SCRYPT

endif

ifdef USE_SCRYPT

endif

ifdef USE_SCRYPT

endif

include "findnonce.h"

include "scrypt.h"

ifdef USE_USBUTILS

ifdef USE_SCRYPT

else

define opt_scrypt (0)

endif

endif /* MINER_H */

define quarter(a,b,c,d) \

undef quarter

pragma unroll

pragma unroll

pragma unroll

pragma unroll

pragma unroll

pragma unroll

pragma unroll

pragma unroll

pragma unroll

pragma unroll

pragma unroll

pragma unroll

pragma unroll

pragma unroll

pragma unroll

pragma unroll

pragma unroll

pragma unroll

pragma unroll

pragma unroll