Merge branch 'nfactor'

Conflicts (resolved):
	Makefile.am
	winbuild/sgminer.vcxproj.filters

Makefile.am

@@ -42,6 +42,7 @@ sgminer_SOURCES += ocl.c ocl.h
 sgminer_SOURCES += findnonce.c findnonce.h
 sgminer_SOURCES += adl.c adl.h adl_functions.h
 sgminer_SOURCES += pool.c pool.h
+sgminer_SOURCES += algorithm.c algorithm.h
 sgminer_SOURCES += scrypt.c scrypt.h
 sgminer_SOURCES += kernel/*.cl
 

algorithm.c

@@ -0,0 +1,35 @@
+/*
+ * Copyright 2014 sgminer developers
+ *
+ * This program is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU General Public License as published by
+ * the Free Software Foundation; either version 3 of the License, or (at
+ * your option) any later version.  See COPYING for more details.
+ */
+
+#include "algorithm.h"
+
+#include <inttypes.h>
+#include <string.h>
+
+void set_algorithm(algorithm_t* algo, const char* newname) {
+    strncpy(algo->name, newname, sizeof(algo->name));
+    algo->name[sizeof(algo->name) - 1] = '\0';
+
+    if ((strcmp(algo->name, "adaptive-n-factor") == 0) ||
+	(strcmp(algo->name, "adaptive-nfactor")  == 0) ||
+	(strcmp(algo->name, "nscrypt")           == 0) ) {
+	set_algorithm_nfactor(algo, 11);
+    } else {
+	set_algorithm_nfactor(algo, 10);
+    }
+
+    return;
+}
+
+void set_algorithm_nfactor(algorithm_t* algo, const uint8_t nfactor) {
+    algo->nfactor = nfactor;
+    algo->n = (1 << nfactor);
+
+    return;
+}

algorithm.h

@@ -0,0 +1,21 @@
+#ifndef ALGORITHM_H
+#define ALGORITHM_H
+
+#include <inttypes.h>
+
+/* Describes the Scrypt parameters and hashing functions used to mine
+ * a specific coin.
+ */
+typedef struct _algorithm_t {
+    char     name[20]; /* Human-readable identifier */
+    uint32_t n;        /* N (CPU/Memory tradeoff parameter) */
+    uint8_t  nfactor;  /* Factor of N above (n = 2^nfactor) */
+} algorithm_t;
+
+/* Set default parameters based on name. */
+void set_algorithm(algorithm_t* algo, const char* name);
+
+/* Set to specific N factor. */
+void set_algorithm_nfactor(algorithm_t* algo, const uint8_t nfactor);
+
+#endif /* ALGORITHM_H */

api.c

@@ -2856,7 +2856,7 @@ static void minecoin(struct io_data *io_data, __maybe_unused SOCKETTYPE c, __may
 	message(io_data, MSG_MINECOIN, 0, NULL, isjson);
 	io_open = io_add(io_data, isjson ? COMSTR JSON_MINECOIN : _MINECOIN COMSTR);
 
-	root = api_add_const(root, "Hash Method", SCRYPTSTR, false);
+	root = api_add_string(root, "Hash Method", algorithm->name, false);
 
 	cg_rlock(&ch_lock);
 	root = api_add_timeval(root, "Current Block Time", &block_timeval, true);

doc/configuration.md

@@ -0,0 +1,36 @@
+# Configuration and command-line options
+
+*Work in progress!*
+
+
+## Config-file and CLI options
+
+### algorithm
+
+Allows choosing between the few mining algorithms for incompatible
+cryptocurrencies.
+
+*Argument:* string
+
+*Default:* `default`
+
+*Supported:*
+
+* `adaptive-nfactor` - Vertcoin-style adaptive N-factor scrypt.
+N-factor defaults to 11.
+* everything else - Litecoin-style static N-factor scrypt.
+
+
+### nfactor
+
+Overrides the default scrypt parameter N, specified as the factor of 2
+(`N = 2^nfactor`).
+
+*Argument:* whole number (>1).
+
+*Default:* depends on `algorithm`; otherwise `10`.
+
+
+## CLI-only options
+
+*TODO*

kernel/alexkarnew.cl

@@ -28,6 +28,36 @@
  * online backup system.
  */
 
+/* N (nfactor), CPU/Memory cost parameter */
+__constant uint N[] = {
+	0x00000001U,  /* never used, padding */
+	0x00000002U,
+	0x00000004U,
+	0x00000008U,
+	0x00000010U,
+	0x00000020U,
+	0x00000040U,
+	0x00000080U,
+	0x00000100U,
+	0x00000200U,
+	0x00000400U,  /* 2^10 == 1024, Litecoin scrypt default */
+	0x00000800U,
+	0x00001000U,
+	0x00002000U,
+	0x00004000U,
+	0x00008000U,
+	0x00010000U,
+	0x00020000U,
+	0x00040000U,
+	0x00080000U,
+	0x00100000U
+};
+
+/* Backwards compatibility, if NFACTOR not defined, default to 10 for scrypt */
+#ifndef NFACTOR
+#define NFACTOR 10
+#endif
+
 __constant uint ES[2] = { 0x00FF00FF, 0xFF00FF00 };
 __constant uint K[] = {
 	0x428a2f98U,
@@ -761,13 +791,13 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 	uint CO=rotl(x,3U);
 	uint CO_tmp=rotl(xSIZE,3U);
 
-	for(uint y=0; y<1024/LOOKUP_GAP; ++y, CO+=CO_tmp)
+	for(uint y=0; y<N[NFACTOR]/LOOKUP_GAP; ++y, CO+=CO_tmp)
 	{
 		uint CO_reg=CO;
 #pragma unroll
 		for(uint z=0; z<zSIZE; ++z, ++CO_reg)
 			lookup[CO_reg] = X[z];
-		for(uint i=0; i<LOOKUP_GAP; ++i) 
+		for(uint i=0; i<LOOKUP_GAP; ++i)
 			salsa(X);
 	}
 
@@ -775,20 +805,20 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 
 #if (LOOKUP_GAP != 1) && (LOOKUP_GAP != 2) && (LOOKUP_GAP != 4) && (LOOKUP_GAP != 8)
 	{
-		uint y = (1024/LOOKUP_GAP);
+		uint y = (N[NFACTOR]/LOOKUP_GAP);
 		CO=CO_tmp+rotl(y*xSIZE,3U);
 #pragma unroll
 		for(uint z=0; z<zSIZE; ++z, ++CO)
 			lookup[CO] = X[z];
-		for(uint i=0; i<1024%LOOKUP_GAP; ++i)
+		for(uint i=0; i<N[NFACTOR]%LOOKUP_GAP; ++i)
 			salsa(X); 
 	}
 #endif
 
-	for (uint i=0; i<1024; ++i) 
+	for (uint i=0; i<N[NFACTOR]; ++i)
 	{
 		uint4 V[8];
-		uint j = X[7].x & K[85];
+		uint j = X[7].x & (N[NFACTOR]-1);
 		uint y = (j/LOOKUP_GAP);
 		uint CO_reg=CO_tmp+rotl(xSIZE*y,3U);
 
@@ -801,7 +831,7 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 			salsa(V);
 #else
 		uint val = j%LOOKUP_GAP;
-		for (uint z=0; z<val; ++z) 
+		for (uint z=0; z<val; ++z)
 			salsa(V);
 #endif
 
@@ -836,7 +866,7 @@ const uint4 midstate0, const uint4 midstate16, const uint target)
 	SHA256(&tstate0, &tstate1, input[0],input[1],input[2],input[3]);
 
 #pragma unroll
-	for (uint i=0; i<4; i++) 
+	for (uint i=0; i<4; i++)
 	{
 		pad0 = tstate0;
 		pad1 = tstate1;

kernel/alexkarold.cl

@@ -28,6 +28,36 @@
  * online backup system.
  */
 
+/* N (nfactor), CPU/Memory cost parameter */
+__constant uint N[] = {
+	0x00000001U,  /* never used, padding */
+	0x00000002U,
+	0x00000004U,
+	0x00000008U,
+	0x00000010U,
+	0x00000020U,
+	0x00000040U,
+	0x00000080U,
+	0x00000100U,
+	0x00000200U,
+	0x00000400U,  /* 2^10 == 1024, Litecoin scrypt default */
+	0x00000800U,
+	0x00001000U,
+	0x00002000U,
+	0x00004000U,
+	0x00008000U,
+	0x00010000U,
+	0x00020000U,
+	0x00040000U,
+	0x00080000U,
+	0x00100000U
+};
+
+/* Backwards compatibility, if NFACTOR not defined, default to 10 for scrypt */
+#ifndef NFACTOR
+#define NFACTOR 10
+#endif
+
 __constant uint ES[2] = { 0x00FF00FF, 0xFF00FF00 };
 __constant uint K[] = {
 	0x428a2f98U,
@@ -761,31 +791,31 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 	uint CO_tmp=xSIZE<<3U;
 	uint CO_tmp2=x<<3U;
 
-	for(uint y=0; y<1024/LOOKUP_GAP; ++y)
+	for(uint y=0; y<N[NFACTOR]/LOOKUP_GAP; ++y)
 	{
 		uint CO=y*CO_tmp+CO_tmp2;
 #pragma unroll
 		for(uint z=0; z<zSIZE; ++z,++CO)
 			lookup[CO] = X[z];
-		for(uint i=0; i<LOOKUP_GAP; ++i) 
+		for(uint i=0; i<LOOKUP_GAP; ++i)
 			salsa(X);
 	}
 	
 #if (LOOKUP_GAP != 1) && (LOOKUP_GAP != 2) && (LOOKUP_GAP != 4) && (LOOKUP_GAP != 8)
 	{
-		uint y = (1024/LOOKUP_GAP);
+		uint y = (N[NFACTOR]/LOOKUP_GAP);
 		uint CO=y*CO_tmp+CO_tmp2;
 #pragma unroll
 		for(uint z=0; z<zSIZE; ++z)
 			lookup[CO] = X[z];
-		for(uint i=0; i<1024%LOOKUP_GAP; ++i)
+		for(uint i=0; i<N[NFACTOR]%LOOKUP_GAP; ++i)
 			salsa(X); 
 	}
 #endif
-	for (uint i=0; i<1024; ++i) 
+	for (uint i=0; i<N[NFACTOR]; ++i)
 	{
 		uint4 V[8];
-		uint j = X[7].x & K[85];
+		uint j = X[7].x & (N[NFACTOR]-1);
 		uint y = (j/LOOKUP_GAP);
 		uint CO=y*CO_tmp+CO_tmp2;
 #pragma unroll
@@ -798,7 +828,7 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 			salsa(V);
 #else
 		uint val = j%LOOKUP_GAP;
-		for (uint z=0; z<val; ++z) 
+		for (uint z=0; z<val; ++z)
 			salsa(V);
 #endif
 
@@ -833,7 +863,7 @@ const uint4 midstate0, const uint4 midstate16, const uint target)
 	SHA256(&tstate0, &tstate1, input[0],input[1],input[2],input[3]);
 
 #pragma unroll
-	for (uint i=0; i<4; i++) 
+	for (uint i=0; i<4; i++)
 	{
 		pad0 = tstate0;
 		pad1 = tstate1;

kernel/ckolivas.cl

@@ -28,6 +28,36 @@
  * online backup system.
  */
 
+/* N (nfactor), CPU/Memory cost parameter */
+__constant uint N[] = {
+	0x00000001U,  /* never used, padding */
+	0x00000002U,
+	0x00000004U,
+	0x00000008U,
+	0x00000010U,
+	0x00000020U,
+	0x00000040U,
+	0x00000080U,
+	0x00000100U,
+	0x00000200U,
+	0x00000400U,  /* 2^10 == 1024, Litecoin scrypt default */
+	0x00000800U,
+	0x00001000U,
+	0x00002000U,
+	0x00004000U,
+	0x00008000U,
+	0x00010000U,
+	0x00020000U,
+	0x00040000U,
+	0x00080000U,
+	0x00100000U
+};
+
+/* Backwards compatibility, if NFACTOR not defined, default to 10 for scrypt */
+#ifndef NFACTOR
+#define NFACTOR 10
+#endif
+
 __constant uint ES[2] = { 0x00FF00FF, 0xFF00FF00 };
 __constant uint K[] = {
 	0x428a2f98U,
@@ -759,32 +789,32 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 {
 	shittify(X);
 	const uint zSIZE = 8;
-	const uint ySIZE = (1024/LOOKUP_GAP+(1024%LOOKUP_GAP>0));
+	const uint ySIZE = (N[NFACTOR]/LOOKUP_GAP+(N[NFACTOR]%LOOKUP_GAP>0));
 	const uint xSIZE = CONCURRENT_THREADS;
 	uint x = get_global_id(0)%xSIZE;
 
-	for(uint y=0; y<1024/LOOKUP_GAP; ++y)
+	for(uint y=0; y<N[NFACTOR]/LOOKUP_GAP; ++y)
 	{
 #pragma unroll
 		for(uint z=0; z<zSIZE; ++z)
 			lookup[CO] = X[z];
-		for(uint i=0; i<LOOKUP_GAP; ++i) 
+		for(uint i=0; i<LOOKUP_GAP; ++i)
 			salsa(X);
 	}
 #if (LOOKUP_GAP != 1) && (LOOKUP_GAP != 2) && (LOOKUP_GAP != 4) && (LOOKUP_GAP != 8)
 	{
-		uint y = (1024/LOOKUP_GAP);
+		uint y = (N[NFACTOR]/LOOKUP_GAP);
 #pragma unroll
 		for(uint z=0; z<zSIZE; ++z)
 			lookup[CO] = X[z];
-		for(uint i=0; i<1024%LOOKUP_GAP; ++i)
+		for(uint i=0; i<N[NFACTOR]%LOOKUP_GAP; ++i)
 			salsa(X); 
 	}
 #endif
-	for (uint i=0; i<1024; ++i) 
+	for (uint i=0; i<N[NFACTOR]; ++i)
 	{
 		uint4 V[8];
-		uint j = X[7].x & K[85];
+		uint j = X[7].x & (N[NFACTOR]-1);
 		uint y = (j/LOOKUP_GAP);
 #pragma unroll
 		for(uint z=0; z<zSIZE; ++z)
@@ -796,7 +826,7 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 			salsa(V);
 #else
 		uint val = j%LOOKUP_GAP;
-		for (uint z=0; z<val; ++z) 
+		for (uint z=0; z<val; ++z)
 			salsa(V);
 #endif
 
@@ -831,7 +861,7 @@ const uint4 midstate0, const uint4 midstate16, const uint target)
 	SHA256(&tstate0, &tstate1, input[0],input[1],input[2],input[3]);
 
 #pragma unroll
-	for (uint i=0; i<4; i++) 
+	for (uint i=0; i<4; i++)
 	{
 		pad0 = tstate0;
 		pad1 = tstate1;

kernel/psw.cl

@@ -29,6 +29,36 @@
  * online backup system.
  */
 
+/* N (nfactor), CPU/Memory cost parameter */
+__constant uint N[] = {
+	0x00000001U,  /* never used, padding */
+	0x00000002U,
+	0x00000004U,
+	0x00000008U,
+	0x00000010U,
+	0x00000020U,
+	0x00000040U,
+	0x00000080U,
+	0x00000100U,
+	0x00000200U,
+	0x00000400U,  /* 2^10 == 1024, Litecoin scrypt default */
+	0x00000800U,
+	0x00001000U,
+	0x00002000U,
+	0x00004000U,
+	0x00008000U,
+	0x00010000U,
+	0x00020000U,
+	0x00040000U,
+	0x00080000U,
+	0x00100000U
+};
+
+/* Backwards compatibility, if NFACTOR not defined, default to 10 for scrypt */
+#ifndef NFACTOR
+#define NFACTOR 10
+#endif
+
 __constant uint ES[2] = { 0x00FF00FF, 0xFF00FF00 };
 __constant uint K[] = {
 	0x428a2f98U,
@@ -698,32 +728,32 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 {
 	shittify(X);
 	const uint zSIZE = 8;
-	const uint ySIZE = (1024/LOOKUP_GAP+(1024%LOOKUP_GAP>0));
+	const uint ySIZE = (N[NFACTOR]/LOOKUP_GAP+(N[NFACTOR]%LOOKUP_GAP>0));
 	const uint xSIZE = CONCURRENT_THREADS;
 	uint x = get_global_id(0)%xSIZE;
 
-	for(uint y=0; y<1024/LOOKUP_GAP; ++y)
+	for(uint y=0; y<N[NFACTOR]/LOOKUP_GAP; ++y)
 	{
 #pragma unroll
 		for(uint z=0; z<zSIZE; ++z)
 			lookup[CO] = X[z];
-		for(uint i=0; i<LOOKUP_GAP; ++i) 
+		for(uint i=0; i<LOOKUP_GAP; ++i)
 			salsa(X);
 	}
 #if (LOOKUP_GAP != 1) && (LOOKUP_GAP != 2) && (LOOKUP_GAP != 4) && (LOOKUP_GAP != 8)
 	{
-		uint y = (1024/LOOKUP_GAP);
+		uint y = (N[NFACTOR]/LOOKUP_GAP);
 #pragma unroll
 		for(uint z=0; z<zSIZE; ++z)
 			lookup[CO] = X[z];
-		for(uint i=0; i<1024%LOOKUP_GAP; ++i)
+		for(uint i=0; i<N[NFACTOR]%LOOKUP_GAP; ++i)
 			salsa(X); 
 	}
 #endif
-	for (uint i=0; i<1024; ++i) 
+	for (uint i=0; i<N[NFACTOR]; ++i)
 	{
 		uint4 V[8];
-		uint j = X[7].x & K[85];
+		uint j = X[7].x & (N[NFACTOR]-1);
 		uint y = (j/LOOKUP_GAP);
 #pragma unroll
 		for(uint z=0; z<zSIZE; ++z)
@@ -735,7 +765,7 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 			salsa(V);
 #else
 		uint val = j%LOOKUP_GAP;
-		for (uint z=0; z<val; ++z) 
+		for (uint z=0; z<val; ++z)
 			salsa(V);
 #endif
 
@@ -770,7 +800,7 @@ const uint4 midstate0, const uint4 midstate16, const uint target)
 	SHA256(&tstate0, &tstate1, input[0],input[1],input[2],input[3]);
 
 #pragma unroll
-	for (uint i=0; i<4; i++) 
+	for (uint i=0; i<4; i++)
 	{
 		pad0 = tstate0;
 		pad1 = tstate1;

kernel/zuikkis.cl

@@ -28,6 +28,36 @@
  * online backup system.
  */
 
+/* N (nfactor), CPU/Memory cost parameter */
+__constant uint N[] = {
+	0x00000001U,  /* never used, padding */
+	0x00000002U,
+	0x00000004U,
+	0x00000008U,
+	0x00000010U,
+	0x00000020U,
+	0x00000040U,
+	0x00000080U,
+	0x00000100U,
+	0x00000200U,
+	0x00000400U,  /* 2^10 == 1024, Litecoin scrypt default */
+	0x00000800U,
+	0x00001000U,
+	0x00002000U,
+	0x00004000U,
+	0x00008000U,
+	0x00010000U,
+	0x00020000U,
+	0x00040000U,
+	0x00080000U,
+	0x00100000U
+};
+
+/* Backwards compatibility, if NFACTOR not defined, default to 10 for scrypt */
+#ifndef NFACTOR
+#define NFACTOR 10
+#endif
+
 __constant uint ES[2] = { 0x00FF00FF, 0xFF00FF00 };
 __constant uint K[] = {
 	0x428a2f98U,
@@ -759,21 +789,21 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 {
 	shittify(X);
 	const uint zSIZE = 8;
-	const uint ySIZE = (1024/LOOKUP_GAP+(1024%LOOKUP_GAP>0));
+	const uint ySIZE = (N[NFACTOR]/LOOKUP_GAP+(N[NFACTOR]%LOOKUP_GAP>0));
 	const uint xSIZE = CONCURRENT_THREADS;
 	uint x = get_global_id(0)%xSIZE;
 
-	for(uint y=0; y<1024/LOOKUP_GAP; ++y)
+	for(uint y=0; y<(N[NFACTOR]/LOOKUP_GAP); ++y)
 	{
 
 		for(uint z=0; z<zSIZE; ++z)
 			lookup[CO] = X[z];
-		for(uint i=0; i<LOOKUP_GAP; ++i) 
+		for(uint i=0; i<LOOKUP_GAP; ++i)
 			salsa(X);
 	}
-	for (uint i=0; i<1024; ++i) 
+	for (uint i=0; i<N[NFACTOR]; ++i)
 	{
-		uint j = X[7].x & K[85];
+		uint j = X[7].x & (N[NFACTOR]-1);
 		uint y = (j/LOOKUP_GAP);
 
 		if (j&1)
@@ -819,15 +849,15 @@ const uint4 midstate0, const uint4 midstate16, const uint target)
 	SHA256(&tstate0, &tstate1, input[0],input[1],input[2],input[3]);
 
 
-	for (uint i=0; i<4; i++) 
+	for (uint i=0; i<4; i++)
 	{
 		pad0 = tstate0;
 		pad1 = tstate1;
-		X[i<<1 ] = ostate0;
+		X[i*2 ] = ostate0;
-		X[(i<<1)+1] = ostate1;
+		X[i*2+1] = ostate1;
 
 		SHA256(&pad0,&pad1, data, (uint4)(i+1,K[84],0,0), (uint4)(0,0,0,0), (uint4)(0,0,0, K[87]));
-		SHA256(X+(i<<1),X+(i<<1)+1, pad0, pad1, (uint4)(K[84], 0U, 0U, 0U), (uint4)(0U, 0U, 0U, K[88]));
+		SHA256(X+i*2,X+i*2+1, pad0, pad1, (uint4)(K[84], 0U, 0U, 0U), (uint4)(0U, 0U, 0U, K[88]));
 	}
 	scrypt_core(X,padcache);
 	SHA256(&tmp0,&tmp1, X[0], X[1], X[2], X[3]);

miner.h

@@ -3,6 +3,8 @@
 
 #include "config.h"
 
+#include "algorithm.h"
+
 #include <stdbool.h>
 #include <stdint.h>
 #include <sys/time.h>
@@ -1020,6 +1022,9 @@ extern int opt_queue;
 extern int opt_scantime;
 extern int opt_expiry;
 
+extern char *opt_algorithm;
+extern algorithm_t *algorithm;
+
 extern cglock_t control_lock;
 extern pthread_mutex_t hash_lock;
 extern pthread_mutex_t console_lock;

ocl.c

@@ -33,6 +33,12 @@
 #include "findnonce.h"
 #include "ocl.h"
 
+/* FIXME: only here for global config vars, replace with configuration.h
+ * or similar as soon as config is in a struct instead of littered all
+ * over the global namespace.
+ */
+#include "miner.h"
+
 int opt_platform_id = -1;
 
 char *file_contents(const char *filename, int *length)
@@ -456,7 +462,8 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
 			break;
 	}
 
-	if (cgpu->vwidth)
+	/* Vectors are hard-set to 1 above. */
+	if (likely(cgpu->vwidth))
 		clState->vwidth = cgpu->vwidth;
 	else {
 		clState->vwidth = preferred_vwidth;
@@ -479,7 +486,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
 	if (!cgpu->opt_tc) {
 		unsigned int sixtyfours;
 
-		sixtyfours =  cgpu->max_alloc / 131072 / 64 - 1;
+		sixtyfours =  cgpu->max_alloc / 131072 / 64 / (algorithm->n/1024) - 1;
 		cgpu->thread_concurrency = sixtyfours * 64;
 		if (cgpu->shaders && cgpu->thread_concurrency > cgpu->shaders) {
 			cgpu->thread_concurrency -= cgpu->thread_concurrency % cgpu->shaders;
@@ -519,7 +526,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
 	if (clState->goffset)
 		strcat(binaryfilename, "g");
 
-	sprintf(numbuf, "lg%utc%u", cgpu->lookup_gap, (unsigned int)cgpu->thread_concurrency);
+	sprintf(numbuf, "lg%utc%unf%u", cgpu->lookup_gap, (unsigned int)cgpu->thread_concurrency, algorithm->nfactor);
 	strcat(binaryfilename, numbuf);
 
 	sprintf(numbuf, "w%d", (int)clState->wsize);
@@ -585,8 +592,8 @@ build:
 	/* create a cl program executable for all the devices specified */
 	char *CompilerOptions = (char *)calloc(1, 256);
 
-	sprintf(CompilerOptions, "-D LOOKUP_GAP=%d -D CONCURRENT_THREADS=%d -D WORKSIZE=%d",
+	sprintf(CompilerOptions, "-D LOOKUP_GAP=%d -D CONCURRENT_THREADS=%d -D WORKSIZE=%d -D NFACTOR=%d",
-			cgpu->lookup_gap, (unsigned int)cgpu->thread_concurrency, (int)clState->wsize);
+			cgpu->lookup_gap, (unsigned int)cgpu->thread_concurrency, (int)clState->wsize, (unsigned int)algorithm->nfactor);
 
 	applog(LOG_DEBUG, "Setting worksize to %d", (int)(clState->wsize));
 	if (clState->vwidth > 1)
@@ -775,7 +782,8 @@ built:
 		return NULL;
 	}
 
-	size_t ipt = (1024 / cgpu->lookup_gap + (1024 % cgpu->lookup_gap > 0));
+	size_t ipt = (algorithm->n / cgpu->lookup_gap +
+		      (algorithm->n % cgpu->lookup_gap > 0));
 	size_t bufsize = 128 * ipt * cgpu->thread_concurrency;
 
 	/* Use the max alloc value which has been rounded to a power of

scrypt.c

@@ -353,10 +353,12 @@ salsa20_8(uint32_t B[16], const uint32_t Bx[16])
 	B[15] += x15;
 }
 
-/* cpu and memory intensive function to transform a 80 byte buffer into a 32 byte output
+/* cpu and memory intensive function to transform a 80 byte buffer into
-   scratchpad size needs to be at least 63 + (128 * r * p) + (256 * r + 64) + (128 * r * N) bytes
+ * a 32 byte output.
+ * scratchpad size needs to be at least (bytes):
+ * 63 + (128 * r * p) + (256 * r + 64) + (128 * r * N)
  */
-static void scrypt_1024_1_1_256_sp(const uint32_t* input, char* scratchpad, uint32_t *ostate)
+static void scrypt_n_1_1_256_sp(const uint32_t* input, char* scratchpad, uint32_t *ostate)
 {
 	uint32_t * V;
 	uint32_t X[32];
@@ -370,7 +372,7 @@ static void scrypt_1024_1_1_256_sp(const uint32_t* input, char* scratchpad, uint
 
 	PBKDF2_SHA256_80_128(input, X);
 
-	for (i = 0; i < 1024; i += 2) {
+	for (i = 0; i < algorithm->n; i += 2) {
 		memcpy(&V[i * 32], X, 128);
 
 		salsa20_8(&X[0], &X[16]);
@@ -381,8 +383,8 @@ static void scrypt_1024_1_1_256_sp(const uint32_t* input, char* scratchpad, uint
 		salsa20_8(&X[0], &X[16]);
 		salsa20_8(&X[16], &X[0]);
 	}
-	for (i = 0; i < 1024; i += 2) {
+	for (i = 0; i < algorithm->n; i += 2) {
-		j = X[16] & 1023;
+		j = X[16] & (algorithm->n-1);
 		p2 = (uint64_t *)(&V[j * 32]);
 		for(k = 0; k < 16; k++)
 			p1[k] ^= p2[k];
@@ -390,7 +392,7 @@ static void scrypt_1024_1_1_256_sp(const uint32_t* input, char* scratchpad, uint
 		salsa20_8(&X[0], &X[16]);
 		salsa20_8(&X[16], &X[0]);
 
-		j = X[16] & 1023;
+		j = X[16] & (algorithm->n-1);
 		p2 = (uint64_t *)(&V[j * 32]);
 		for(k = 0; k < 16; k++)
 			p1[k] ^= p2[k];
@@ -402,26 +404,26 @@ static void scrypt_1024_1_1_256_sp(const uint32_t* input, char* scratchpad, uint
 	PBKDF2_SHA256_80_128_32(input, X, ostate);
 }
 
-/* 131583 rounded up to 4 byte alignment */
-#define SCRATCHBUF_SIZE	(131584)
-
 void scrypt_regenhash(struct work *work)
 {
 	uint32_t data[20];
 	char *scratchbuf;
 	uint32_t *nonce = (uint32_t *)(work->data + 76);
 	uint32_t *ohash = (uint32_t *)(work->hash);
-
+ 	
 	be32enc_vect(data, (const uint32_t *)work->data, 19);
 	data[19] = htobe32(*nonce);
-	scratchbuf = (char *)alloca(SCRATCHBUF_SIZE);
+
-	scrypt_1024_1_1_256_sp(data, scratchbuf, ohash);
+	scratchbuf = (char *)alloca(algorithm->n * 128 + 512);
+	scrypt_n_1_1_256_sp(data, scratchbuf, ohash);
 	flip32(ohash, ohash);
 }
 
 static const uint32_t diff1targ = 0x0000ffff;
 
 /* Used externally as confirmation of correct OCL code */
+/* FIXME: find reference in git blame and see why it was present, remove if obsolete */
+/*
 int scrypt_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t nonce)
 {
 	uint32_t tmp_hash7, Htarg = le32toh(((const uint32_t *)ptarget)[7]);
@@ -431,7 +433,7 @@ int scrypt_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t non
 	be32enc_vect(data, (const uint32_t *)pdata, 19);
 	data[19] = htobe32(nonce);
 	scratchbuf = (char *)alloca(SCRATCHBUF_SIZE);
-	scrypt_1024_1_1_256_sp(data, scratchbuf, ohash);
+	scrypt_n_1_1_256_sp(data, scratchbuf, ohash, algorithm->n);
 	tmp_hash7 = be32toh(ohash[7]);
 
 	applog(LOG_DEBUG, "htarget %08lx diff1 %08lx hash %08lx",
@@ -470,7 +472,7 @@ bool scanhash_scrypt(struct thr_info *thr, const unsigned char __maybe_unused *p
 
 		*nonce = ++n;
 		data[19] = htobe32(n);
-		scrypt_1024_1_1_256_sp(data, scratchbuf, ostate);
+		scrypt_n_1_1_256_sp(data, scratchbuf, ostate, algorithm->n);
 		tmp_hash7 = be32toh(ostate[7]);
 
 		if (unlikely(tmp_hash7 <= Htarg)) {
@@ -489,3 +491,4 @@ bool scanhash_scrypt(struct thr_info *thr, const unsigned char __maybe_unused *p
 	free(scratchbuf);;
 	return ret;
 }
+*/

scrypt.h

@@ -3,8 +3,8 @@
 
 #include "miner.h"
 
-extern int scrypt_test(unsigned char *pdata, const unsigned char *ptarget,
+/* extern int scrypt_test(unsigned char *pdata, const unsigned char *ptarget, */
-			uint32_t nonce);
+/* 			uint32_t nonce); */
 extern void scrypt_regenhash(struct work *work);
 
 #endif /* SCRYPT_H */

sgminer.c

@@ -54,6 +54,8 @@ char *curly = ":D";
 #include "adl.h"
 #include "driver-opencl.h"
 #include "bench_block.h"
+
+#include "algorithm.h"
 #include "scrypt.h"
 #include "pool.h"
 
@@ -95,6 +97,10 @@ int opt_log_interval = 5;
 int opt_queue = 1;
 int opt_scantime = 7;
 int opt_expiry = 28;
+
+char *opt_algorithm;
+algorithm_t *algorithm;
+
 static const bool opt_time = true;
 unsigned long long global_hashrate;
 unsigned long global_quota_gcd = 1;
@@ -1007,6 +1013,23 @@ static void load_temp_cutoffs()
 	}
 }
 
+static char *set_algo(const char *arg)
+{
+	set_algorithm(algorithm, arg);
+	applog(LOG_INFO, "Set algorithm to %s", algorithm->name);
+
+	return NULL;
+}
+
+static char *set_nfactor(const char *arg)
+{
+	set_algorithm_nfactor(algorithm, (uint8_t)atoi(arg));
+	applog(LOG_INFO, "Set algorithm N-factor to %d (N to %d)",
+	       algorithm->nfactor, algorithm->n);
+
+	return NULL;
+}
+
 static char *set_api_allow(const char *arg)
 {
 	opt_set_charp(arg, &opt_api_allow);
@@ -1056,6 +1079,9 @@ static char *set_null(const char __maybe_unused *arg)
 
 /* These options are available from config file or commandline */
 static struct opt_table opt_config_table[] = {
+	OPT_WITH_ARG("--algorithm",
+		     set_algo, NULL, NULL,
+		     "Set mining algorithm and most common defaults, default: static"),
 	OPT_WITH_ARG("--api-allow",
 		     set_api_allow, NULL, NULL,
 		     "Allow API access only to the given list of [G:]IP[/Prefix] addresses[/subnets]"),
@@ -1216,13 +1242,16 @@ static struct opt_table opt_config_table[] = {
 	OPT_WITHOUT_ARG("--net-delay",
 			opt_set_bool, &opt_delaynet,
 			"Impose small delays in networking to not overload slow routers"),
+	OPT_WITH_ARG("--nfactor",
+		     set_nfactor, NULL, NULL,
+		     "Override default scrypt N-factor parameter."),
 #ifdef HAVE_ADL
 	OPT_WITHOUT_ARG("--no-adl",
 			opt_set_bool, &opt_noadl,
 			"Disable the ATI display library used for monitoring and setting GPU parameters"),
 #else
 	OPT_WITHOUT_ARG("--no-adl",
-			opt_set_bool, &opt_noadl,opt_hidden),
+			opt_set_bool, &opt_noadl, opt_hidden),
 #endif
 	OPT_WITHOUT_ARG("--no-pool-disable",
 			opt_set_invbool, &opt_disable_pool,
@@ -7767,7 +7796,7 @@ int main(int argc, char *argv[])
 	int i, j;
 	char *s;
 
-	/* This dangerous functions tramples random dynamically allocated
+	/* This dangerous function tramples random dynamically allocated
 	 * variables so do it before anything at all */
 	if (unlikely(curl_global_init(CURL_GLOBAL_ALL)))
 		quit(1, "Failed to curl_global_init");
@@ -7845,6 +7874,10 @@ int main(int argc, char *argv[])
 	strcat(sgminer_path, "\\");
 #endif
 
+	/* Default algorithm specified in algorithm.c ATM */
+	algorithm = (algorithm_t *)alloca(sizeof(algorithm_t));
+	set_algorithm(algorithm, "scrypt");
+
 	devcursor = 8;
 	logstart = devcursor + 1;
 	logcursor = logstart + 1;

winbuild/dist/include/inttypes.h

@@ -0,0 +1,305 @@
+// ISO C9x  compliant inttypes.h for Microsoft Visual Studio
+// Based on ISO/IEC 9899:TC2 Committee draft (May 6, 2005) WG14/N1124 
+// 
+//  Copyright (c) 2006 Alexander Chemeris
+// 
+// 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.
+// 
+//   3. The name of the author may be used to endorse or promote products
+//      derived from this software without specific prior written permission.
+// 
+// THIS SOFTWARE IS PROVIDED BY THE AUTHOR ``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 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.
+// 
+///////////////////////////////////////////////////////////////////////////////
+
+#ifndef _MSC_VER // [
+#error "Use this header only with Microsoft Visual C++ compilers!"
+#endif // _MSC_VER ]
+
+#ifndef _MSC_INTTYPES_H_ // [
+#define _MSC_INTTYPES_H_
+
+#if _MSC_VER > 1000
+#pragma once
+#endif
+
+#include "stdint.h"
+
+// 7.8 Format conversion of integer types
+
+typedef struct {
+   intmax_t quot;
+   intmax_t rem;
+} imaxdiv_t;
+
+// 7.8.1 Macros for format specifiers
+
+#if !defined(__cplusplus) || defined(__STDC_FORMAT_MACROS) // [   See footnote 185 at page 198
+
+// The fprintf macros for signed integers are:
+#define PRId8       "d"
+#define PRIi8       "i"
+#define PRIdLEAST8  "d"
+#define PRIiLEAST8  "i"
+#define PRIdFAST8   "d"
+#define PRIiFAST8   "i"
+
+#define PRId16       "hd"
+#define PRIi16       "hi"
+#define PRIdLEAST16  "hd"
+#define PRIiLEAST16  "hi"
+#define PRIdFAST16   "hd"
+#define PRIiFAST16   "hi"
+
+#define PRId32       "I32d"
+#define PRIi32       "I32i"
+#define PRIdLEAST32  "I32d"
+#define PRIiLEAST32  "I32i"
+#define PRIdFAST32   "I32d"
+#define PRIiFAST32   "I32i"
+
+#define PRId64       "I64d"
+#define PRIi64       "I64i"
+#define PRIdLEAST64  "I64d"
+#define PRIiLEAST64  "I64i"
+#define PRIdFAST64   "I64d"
+#define PRIiFAST64   "I64i"
+
+#define PRIdMAX     "I64d"
+#define PRIiMAX     "I64i"
+
+#define PRIdPTR     "Id"
+#define PRIiPTR     "Ii"
+
+// The fprintf macros for unsigned integers are:
+#define PRIo8       "o"
+#define PRIu8       "u"
+#define PRIx8       "x"
+#define PRIX8       "X"
+#define PRIoLEAST8  "o"
+#define PRIuLEAST8  "u"
+#define PRIxLEAST8  "x"
+#define PRIXLEAST8  "X"
+#define PRIoFAST8   "o"
+#define PRIuFAST8   "u"
+#define PRIxFAST8   "x"
+#define PRIXFAST8   "X"
+
+#define PRIo16       "ho"
+#define PRIu16       "hu"
+#define PRIx16       "hx"
+#define PRIX16       "hX"
+#define PRIoLEAST16  "ho"
+#define PRIuLEAST16  "hu"
+#define PRIxLEAST16  "hx"
+#define PRIXLEAST16  "hX"
+#define PRIoFAST16   "ho"
+#define PRIuFAST16   "hu"
+#define PRIxFAST16   "hx"
+#define PRIXFAST16   "hX"
+
+#define PRIo32       "I32o"
+#define PRIu32       "I32u"
+#define PRIx32       "I32x"
+#define PRIX32       "I32X"
+#define PRIoLEAST32  "I32o"
+#define PRIuLEAST32  "I32u"
+#define PRIxLEAST32  "I32x"
+#define PRIXLEAST32  "I32X"
+#define PRIoFAST32   "I32o"
+#define PRIuFAST32   "I32u"
+#define PRIxFAST32   "I32x"
+#define PRIXFAST32   "I32X"
+
+#define PRIo64       "I64o"
+#define PRIu64       "I64u"
+#define PRIx64       "I64x"
+#define PRIX64       "I64X"
+#define PRIoLEAST64  "I64o"
+#define PRIuLEAST64  "I64u"
+#define PRIxLEAST64  "I64x"
+#define PRIXLEAST64  "I64X"
+#define PRIoFAST64   "I64o"
+#define PRIuFAST64   "I64u"
+#define PRIxFAST64   "I64x"
+#define PRIXFAST64   "I64X"
+
+#define PRIoMAX     "I64o"
+#define PRIuMAX     "I64u"
+#define PRIxMAX     "I64x"
+#define PRIXMAX     "I64X"
+
+#define PRIoPTR     "Io"
+#define PRIuPTR     "Iu"
+#define PRIxPTR     "Ix"
+#define PRIXPTR     "IX"
+
+// The fscanf macros for signed integers are:
+#define SCNd8       "d"
+#define SCNi8       "i"
+#define SCNdLEAST8  "d"
+#define SCNiLEAST8  "i"
+#define SCNdFAST8   "d"
+#define SCNiFAST8   "i"
+
+#define SCNd16       "hd"
+#define SCNi16       "hi"
+#define SCNdLEAST16  "hd"
+#define SCNiLEAST16  "hi"
+#define SCNdFAST16   "hd"
+#define SCNiFAST16   "hi"
+
+#define SCNd32       "ld"
+#define SCNi32       "li"
+#define SCNdLEAST32  "ld"
+#define SCNiLEAST32  "li"
+#define SCNdFAST32   "ld"
+#define SCNiFAST32   "li"
+
+#define SCNd64       "I64d"
+#define SCNi64       "I64i"
+#define SCNdLEAST64  "I64d"
+#define SCNiLEAST64  "I64i"
+#define SCNdFAST64   "I64d"
+#define SCNiFAST64   "I64i"
+
+#define SCNdMAX     "I64d"
+#define SCNiMAX     "I64i"
+
+#ifdef _WIN64 // [
+#  define SCNdPTR     "I64d"
+#  define SCNiPTR     "I64i"
+#else  // _WIN64 ][
+#  define SCNdPTR     "ld"
+#  define SCNiPTR     "li"
+#endif  // _WIN64 ]
+
+// The fscanf macros for unsigned integers are:
+#define SCNo8       "o"
+#define SCNu8       "u"
+#define SCNx8       "x"
+#define SCNX8       "X"
+#define SCNoLEAST8  "o"
+#define SCNuLEAST8  "u"
+#define SCNxLEAST8  "x"
+#define SCNXLEAST8  "X"
+#define SCNoFAST8   "o"
+#define SCNuFAST8   "u"
+#define SCNxFAST8   "x"
+#define SCNXFAST8   "X"
+
+#define SCNo16       "ho"
+#define SCNu16       "hu"
+#define SCNx16       "hx"
+#define SCNX16       "hX"
+#define SCNoLEAST16  "ho"
+#define SCNuLEAST16  "hu"
+#define SCNxLEAST16  "hx"
+#define SCNXLEAST16  "hX"
+#define SCNoFAST16   "ho"
+#define SCNuFAST16   "hu"
+#define SCNxFAST16   "hx"
+#define SCNXFAST16   "hX"
+
+#define SCNo32       "lo"
+#define SCNu32       "lu"
+#define SCNx32       "lx"
+#define SCNX32       "lX"
+#define SCNoLEAST32  "lo"
+#define SCNuLEAST32  "lu"
+#define SCNxLEAST32  "lx"
+#define SCNXLEAST32  "lX"
+#define SCNoFAST32   "lo"
+#define SCNuFAST32   "lu"
+#define SCNxFAST32   "lx"
+#define SCNXFAST32   "lX"
+
+#define SCNo64       "I64o"
+#define SCNu64       "I64u"
+#define SCNx64       "I64x"
+#define SCNX64       "I64X"
+#define SCNoLEAST64  "I64o"
+#define SCNuLEAST64  "I64u"
+#define SCNxLEAST64  "I64x"
+#define SCNXLEAST64  "I64X"
+#define SCNoFAST64   "I64o"
+#define SCNuFAST64   "I64u"
+#define SCNxFAST64   "I64x"
+#define SCNXFAST64   "I64X"
+
+#define SCNoMAX     "I64o"
+#define SCNuMAX     "I64u"
+#define SCNxMAX     "I64x"
+#define SCNXMAX     "I64X"
+
+#ifdef _WIN64 // [
+#  define SCNoPTR     "I64o"
+#  define SCNuPTR     "I64u"
+#  define SCNxPTR     "I64x"
+#  define SCNXPTR     "I64X"
+#else  // _WIN64 ][
+#  define SCNoPTR     "lo"
+#  define SCNuPTR     "lu"
+#  define SCNxPTR     "lx"
+#  define SCNXPTR     "lX"
+#endif  // _WIN64 ]
+
+#endif // __STDC_FORMAT_MACROS ]
+
+// 7.8.2 Functions for greatest-width integer types
+
+// 7.8.2.1 The imaxabs function
+#define imaxabs _abs64
+
+// 7.8.2.2 The imaxdiv function
+
+// This is modified version of div() function from Microsoft's div.c found
+// in %MSVC.NET%\crt\src\div.c
+#ifdef STATIC_IMAXDIV // [
+static
+#else // STATIC_IMAXDIV ][
+_inline
+#endif // STATIC_IMAXDIV ]
+imaxdiv_t __cdecl imaxdiv(intmax_t numer, intmax_t denom)
+{
+   imaxdiv_t result;
+
+   result.quot = numer / denom;
+   result.rem = numer % denom;
+
+   if (numer < 0 && result.rem > 0) {
+      // did division wrong; must fix up
+      ++result.quot;
+      result.rem -= denom;
+   }
+
+   return result;
+}
+
+// 7.8.2.3 The strtoimax and strtoumax functions
+#define strtoimax _strtoi64
+#define strtoumax _strtoui64
+
+// 7.8.2.4 The wcstoimax and wcstoumax functions
+#define wcstoimax _wcstoi64
+#define wcstoumax _wcstoui64
+
+
+#endif // _MSC_INTTYPES_H_ ]

winbuild/sgminer.vcxproj

@@ -241,6 +241,7 @@ exit 0</Command>
     <ClCompile Include="..\adl.c">
       <WholeProgramOptimization Condition="'$(Configuration)|$(Platform)'=='Release|x64'">true</WholeProgramOptimization>
     </ClCompile>
+    <ClCompile Include="..\algorithm.c" />
     <ClCompile Include="..\api.c" />
     <ClCompile Include="..\ccan\opt\helpers.c" />
     <ClCompile Include="..\ccan\opt\opt.c" />
@@ -269,6 +270,7 @@ exit 0</Command>
   </ItemGroup>
   <ItemGroup>
     <ClInclude Include="..\adl.h" />
+    <ClInclude Include="..\algorithm.h" />
     <ClInclude Include="..\arg-nonnull.h" />
     <ClInclude Include="..\bench_block.h" />
     <ClInclude Include="..\c++defs.h" />

winbuild/sgminer.vcxproj.filters

@@ -1,4 +1,4 @@
-<?xml version="1.0" encoding="utf-8"?>
+<?xml version="1.0" encoding="utf-8"?>
 <Project ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
   <ItemGroup>
     <Filter Include="Source Files">
@@ -95,6 +95,9 @@
     <ClCompile Include="..\pool.c">
       <Filter>Source Files</Filter>
     </ClCompile>
+    <ClCompile Include="..\algorithm.c">
+      <Filter>Source Files</Filter>
+    </ClCompile>
   </ItemGroup>
   <ItemGroup>
     <ClInclude Include="..\adl.h">
@@ -169,8 +172,11 @@
     <ClInclude Include="..\pool.h">
       <Filter>Header Files</Filter>
     </ClInclude>
+    <ClInclude Include="..\algorithm.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
   </ItemGroup>
   <ItemGroup>
     <None Include="README.txt" />
   </ItemGroup>
-</Project>
+</Project>