From 73bb1504e999b747692661b2d3cd5f2e9c015f63 Mon Sep 17 00:00:00 2001 From: Teemu Suikki Date: Sun, 23 Feb 2014 04:44:17 +0200 Subject: [PATCH 01/21] Scrypt-nfactor support! Added new configuration parameter "nfactor", which defaults to 10 (normal scrypt). Use 11 for vertcoin. Kernels modified accordingly. --- kernel/alexkarnew.cl | 15 ++++++++++----- kernel/alexkarold.cl | 15 ++++++++++----- kernel/ckolivas.cl | 17 +++++++++++------ kernel/psw.cl | 17 +++++++++++------ kernel/zuikkis.cl | 19 ++++++++++++------- miner.h | 1 + ocl.c | 11 ++++++----- scrypt.c | 22 +++++++++++++--------- sgminer.c | 4 ++++ 9 files changed, 78 insertions(+), 43 deletions(-) diff --git a/kernel/alexkarnew.cl b/kernel/alexkarnew.cl index 757e8114..43486c3e 100644 --- a/kernel/alexkarnew.cl +++ b/kernel/alexkarnew.cl @@ -28,6 +28,11 @@ * online backup system. */ +/* Backwards compatibility, if NFACTOR not defined, default to 1024 scrypt */ +#ifndef NFACTOR +#define NFACTOR 1024 +#endif + __constant uint ES[2] = { 0x00FF00FF, 0xFF00FF00 }; __constant uint K[] = { 0x428a2f98U, @@ -761,7 +766,7 @@ 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; y0)); + const uint ySIZE = (NFACTOR/LOOKUP_GAP+(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; y0)); + const uint ySIZE = (NFACTOR/LOOKUP_GAP+(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; y0)); + const uint ySIZE = (NFACTOR/LOOKUP_GAP+(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<(NFACTOR/LOOKUP_GAP); ++y) { for(uint z=0; zopt_tc) { unsigned int sixtyfours; - sixtyfours = cgpu->max_alloc / 131072 / 64 - 1; + sixtyfours = cgpu->max_alloc / 131072 / 64 / (nfactor/1024)- 1; cgpu->thread_concurrency = sixtyfours * 64; if (cgpu->shaders && cgpu->thread_concurrency > cgpu->shaders) { cgpu->thread_concurrency -= cgpu->thread_concurrency % cgpu->shaders; @@ -521,7 +522,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%un%u", cgpu->lookup_gap, (unsigned int)cgpu->thread_concurrency,opt_nfactor); strcat(binaryfilename, numbuf); sprintf(numbuf, "w%d", (int)clState->wsize); @@ -587,8 +588,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", - cgpu->lookup_gap, (unsigned int)cgpu->thread_concurrency, (int)clState->wsize); + 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,(unsigned int)nfactor); applog(LOG_DEBUG, "Setting worksize to %d", (int)(clState->wsize)); if (clState->vwidth > 1) @@ -777,7 +778,7 @@ built: return NULL; } - size_t ipt = (1024 / cgpu->lookup_gap + (1024 % cgpu->lookup_gap > 0)); + size_t ipt = (nfactor / cgpu->lookup_gap + (nfactor % 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 diff --git a/scrypt.c b/scrypt.c index d135a72d..b2cc005d 100644 --- a/scrypt.c +++ b/scrypt.c @@ -356,7 +356,7 @@ salsa20_8(uint32_t B[16], const uint32_t Bx[16]) /* cpu and memory intensive function to transform a 80 byte buffer into a 32 byte output scratchpad size needs to be at least 63 + (128 * r * p) + (256 * r + 64) + (128 * r * N) bytes */ -static void scrypt_1024_1_1_256_sp(const uint32_t* input, char* scratchpad, uint32_t *ostate) +static void scrypt_1024_1_1_256_sp(const uint32_t* input, char* scratchpad, uint32_t *ostate, const cl_uint n) { uint32_t * V; uint32_t X[32]; @@ -370,7 +370,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 < n; i += 2) { memcpy(&V[i * 32], X, 128); salsa20_8(&X[0], &X[16]); @@ -381,8 +381,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) { - j = X[16] & 1023; + for (i = 0; i < n; i += 2) { + j = X[16] & (n-1); p2 = (uint64_t *)(&V[j * 32]); for(k = 0; k < 16; k++) p1[k] ^= p2[k]; @@ -390,7 +390,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] & (n-1); p2 = (uint64_t *)(&V[j * 32]); for(k = 0; k < 16; k++) p1[k] ^= p2[k]; @@ -403,7 +403,8 @@ static void scrypt_1024_1_1_256_sp(const uint32_t* input, char* scratchpad, uint } /* 131583 rounded up to 4 byte alignment */ -#define SCRATCHBUF_SIZE (131584) +//#define SCRATCHBUF_SIZE (131584) +//#define SCRATCHBUF_SIZE (262207) void scrypt_regenhash(struct work *work) { @@ -411,17 +412,19 @@ void scrypt_regenhash(struct work *work) 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 = alloca(SCRATCHBUF_SIZE); + scratchbuf = (char *)alloca((1 << opt_nfactor) * 128 + 512); + scrypt_1024_1_1_256_sp(data, scratchbuf, ohash, (1 << opt_nfactor)); flip32(ohash, ohash); } static const uint32_t diff1targ = 0x0000ffff; /* Used externally as confirmation of correct OCL code */ +/* int scrypt_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t nonce) { uint32_t tmp_hash7, Htarg = le32toh(((const uint32_t *)ptarget)[7]); @@ -489,3 +492,4 @@ bool scanhash_scrypt(struct thr_info *thr, const unsigned char __maybe_unused *p free(scratchbuf);; return ret; } +*/ diff --git a/sgminer.c b/sgminer.c index 98e4051c..f16026e4 100644 --- a/sgminer.c +++ b/sgminer.c @@ -92,6 +92,7 @@ int opt_log_interval = 5; int opt_queue = 1; int opt_scantime = 7; int opt_expiry = 28; +int opt_nfactor = 11; static const bool opt_time = true; unsigned long long global_hashrate; unsigned long global_quota_gcd = 1; @@ -1105,6 +1106,9 @@ static struct opt_table opt_config_table[] = { opt_set_bool, &opt_compact, "Use compact display without per device statistics"), #endif + OPT_WITH_ARG("--nfactor", + set_int_0_to_9999, opt_show_intval, &opt_nfactor, + "Set scrypt nfactor, default: 10. Currently use 11 for vertcoin!"), OPT_WITHOUT_ARG("--debug|-D", enable_debug, &opt_debug, "Enable debug output"), From 2b9a588e25e86dcf5b30c11229084a7f7c6711f8 Mon Sep 17 00:00:00 2001 From: Zuikkis Date: Sun, 23 Feb 2014 18:20:57 +0200 Subject: [PATCH 02/21] Update sgminer.c & scrypt.c Default opt_nfactor was 11 instead of 10, my typo.. :) Sorry. scrypt_1024_1_1_256_sp renamed to scrypt_n_1_1_256_sp --- scrypt.c | 4 ++-- sgminer.c | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/scrypt.c b/scrypt.c index b2cc005d..acbd9828 100644 --- a/scrypt.c +++ b/scrypt.c @@ -356,7 +356,7 @@ salsa20_8(uint32_t B[16], const uint32_t Bx[16]) /* cpu and memory intensive function to transform a 80 byte buffer into a 32 byte output scratchpad size needs to be at least 63 + (128 * r * p) + (256 * r + 64) + (128 * r * N) bytes */ -static void scrypt_1024_1_1_256_sp(const uint32_t* input, char* scratchpad, uint32_t *ostate, const cl_uint n) +static void scrypt_n_1_1_256_sp(const uint32_t* input, char* scratchpad, uint32_t *ostate, const cl_uint n) { uint32_t * V; uint32_t X[32]; @@ -417,7 +417,7 @@ void scrypt_regenhash(struct work *work) data[19] = htobe32(*nonce); //scratchbuf = alloca(SCRATCHBUF_SIZE); scratchbuf = (char *)alloca((1 << opt_nfactor) * 128 + 512); - scrypt_1024_1_1_256_sp(data, scratchbuf, ohash, (1 << opt_nfactor)); + scrypt_n_1_1_256_sp(data, scratchbuf, ohash, (1 << opt_nfactor)); flip32(ohash, ohash); } diff --git a/sgminer.c b/sgminer.c index f16026e4..e3301337 100644 --- a/sgminer.c +++ b/sgminer.c @@ -92,7 +92,7 @@ int opt_log_interval = 5; int opt_queue = 1; int opt_scantime = 7; int opt_expiry = 28; -int opt_nfactor = 11; +int opt_nfactor = 10; static const bool opt_time = true; unsigned long long global_hashrate; unsigned long global_quota_gcd = 1; From ab9fcb1a14ea2240837b122f3dacef50633cf815 Mon Sep 17 00:00:00 2001 From: Noel Maersk Date: Mon, 24 Feb 2014 16:18:12 +0200 Subject: [PATCH 03/21] ocl: use same type convention for nfactor (`cl_uint` instead of `int`). --- ocl.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ocl.c b/ocl.c index f128636e..98a5a1fb 100644 --- a/ocl.c +++ b/ocl.c @@ -225,7 +225,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize) cl_uint numPlatforms; cl_uint numDevices; cl_int status; - int nfactor = (1< Date: Mon, 24 Feb 2014 16:20:11 +0200 Subject: [PATCH 04/21] ocl: vectors are hard-set to 1, add appropriate "optimisation". --- ocl.c | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/ocl.c b/ocl.c index 98a5a1fb..2bb52c00 100644 --- a/ocl.c +++ b/ocl.c @@ -459,7 +459,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; From 5b42d38f33bc98279315ebeac7a5066017217701 Mon Sep 17 00:00:00 2001 From: Noel Maersk Date: Mon, 24 Feb 2014 16:22:03 +0200 Subject: [PATCH 05/21] ocl/misc: spacing. --- ocl.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ocl.c b/ocl.c index 2bb52c00..9e0a18a8 100644 --- a/ocl.c +++ b/ocl.c @@ -483,7 +483,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize) if (!cgpu->opt_tc) { unsigned int sixtyfours; - sixtyfours = cgpu->max_alloc / 131072 / 64 / (nfactor/1024)- 1; + sixtyfours = cgpu->max_alloc / 131072 / 64 / (nfactor/1024) - 1; cgpu->thread_concurrency = sixtyfours * 64; if (cgpu->shaders && cgpu->thread_concurrency > cgpu->shaders) { cgpu->thread_concurrency -= cgpu->thread_concurrency % cgpu->shaders; From 4b3b2ef5a6c59cffdad0dc4d8dd490266ea1ac18 Mon Sep 17 00:00:00 2001 From: Noel Maersk Date: Mon, 24 Feb 2014 16:24:00 +0200 Subject: [PATCH 06/21] ocl: use 'nf' instead of 'n' in compiled kernel binary name. It is used as a factor number here (e.g. 10, not 1024). --- ocl.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ocl.c b/ocl.c index 9e0a18a8..3f28b765 100644 --- a/ocl.c +++ b/ocl.c @@ -523,7 +523,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize) if (clState->goffset) strcat(binaryfilename, "g"); - sprintf(numbuf, "lg%utc%un%u", cgpu->lookup_gap, (unsigned int)cgpu->thread_concurrency,opt_nfactor); + sprintf(numbuf, "lg%utc%unf%u", cgpu->lookup_gap, (unsigned int)cgpu->thread_concurrency,opt_nfactor); strcat(binaryfilename, numbuf); sprintf(numbuf, "w%d", (int)clState->wsize); From 2135777cac4b1522112f5b097674490150a25dd2 Mon Sep 17 00:00:00 2001 From: Noel Maersk Date: Tue, 25 Feb 2014 06:25:48 +0200 Subject: [PATCH 07/21] scrypt: remove commented-out magic numbers. --- scrypt.c | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/scrypt.c b/scrypt.c index acbd9828..506b7daf 100644 --- a/scrypt.c +++ b/scrypt.c @@ -402,10 +402,6 @@ static void scrypt_n_1_1_256_sp(const uint32_t* input, char* scratchpad, uint32_ PBKDF2_SHA256_80_128_32(input, X, ostate); } -/* 131583 rounded up to 4 byte alignment */ -//#define SCRATCHBUF_SIZE (131584) -//#define SCRATCHBUF_SIZE (262207) - void scrypt_regenhash(struct work *work) { uint32_t data[20]; @@ -415,7 +411,7 @@ void scrypt_regenhash(struct work *work) be32enc_vect(data, (const uint32_t *)work->data, 19); data[19] = htobe32(*nonce); - //scratchbuf = alloca(SCRATCHBUF_SIZE); + scratchbuf = (char *)alloca((1 << opt_nfactor) * 128 + 512); scrypt_n_1_1_256_sp(data, scratchbuf, ohash, (1 << opt_nfactor)); flip32(ohash, ohash); From c4c85ca71c4251311aeab93e4fdcec71c62d06b4 Mon Sep 17 00:00:00 2001 From: Noel Maersk Date: Tue, 25 Feb 2014 06:32:06 +0200 Subject: [PATCH 08/21] misc: add FIXME for scrypt_test() and scanhash_scrypt(). --- scrypt.c | 5 +++-- scrypt.h | 4 ++-- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/scrypt.c b/scrypt.c index 506b7daf..3b2e74de 100644 --- a/scrypt.c +++ b/scrypt.c @@ -420,6 +420,7 @@ void scrypt_regenhash(struct work *work) 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) { @@ -430,7 +431,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, (1 << opt_nfactor)); tmp_hash7 = be32toh(ohash[7]); applog(LOG_DEBUG, "htarget %08lx diff1 %08lx hash %08lx", @@ -469,7 +470,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, (1 << opt_nfactor)); tmp_hash7 = be32toh(ostate[7]); if (unlikely(tmp_hash7 <= Htarg)) { diff --git a/scrypt.h b/scrypt.h index b08edfaf..d967f541 100644 --- a/scrypt.h +++ b/scrypt.h @@ -3,8 +3,8 @@ #include "miner.h" -extern int scrypt_test(unsigned char *pdata, const unsigned char *ptarget, - uint32_t nonce); +/* extern int scrypt_test(unsigned char *pdata, const unsigned char *ptarget, */ +/* uint32_t nonce); */ extern void scrypt_regenhash(struct work *work); #endif /* SCRYPT_H */ From 3afaaf1bfa21aba1bffd913f654edf43c88527ab Mon Sep 17 00:00:00 2001 From: Noel Maersk Date: Wed, 26 Feb 2014 03:45:14 +0200 Subject: [PATCH 09/21] ocl: rename nfactor to N where appropriate. N = 2^nfactor, so `nfactor` was a bad name choice here. --- ocl.c | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/ocl.c b/ocl.c index 3f28b765..de6ea650 100644 --- a/ocl.c +++ b/ocl.c @@ -225,7 +225,9 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize) cl_uint numPlatforms; cl_uint numDevices; cl_int status; - cl_uint nfactor = (1<opt_tc) { unsigned int sixtyfours; - sixtyfours = cgpu->max_alloc / 131072 / 64 / (nfactor/1024) - 1; + sixtyfours = cgpu->max_alloc / 131072 / 64 / (N/1024) - 1; cgpu->thread_concurrency = sixtyfours * 64; if (cgpu->shaders && cgpu->thread_concurrency > cgpu->shaders) { cgpu->thread_concurrency -= cgpu->thread_concurrency % cgpu->shaders; @@ -779,7 +781,7 @@ built: return NULL; } - size_t ipt = (nfactor / cgpu->lookup_gap + (nfactor % cgpu->lookup_gap > 0)); + size_t ipt = (N / cgpu->lookup_gap + (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 From e18bcc25822c3a09053140e84819e963d862866f Mon Sep 17 00:00:00 2001 From: Noel Maersk Date: Wed, 26 Feb 2014 04:30:46 +0200 Subject: [PATCH 10/21] kernel: use N-factor table instead of passing parameter N directly. --- kernel/alexkarnew.cl | 39 ++++++++++++++++++++++++++++++++------- kernel/alexkarold.cl | 39 ++++++++++++++++++++++++++++++++------- kernel/ckolivas.cl | 41 +++++++++++++++++++++++++++++++++-------- kernel/psw.cl | 41 +++++++++++++++++++++++++++++++++-------- kernel/zuikkis.cl | 37 +++++++++++++++++++++++++++++++------ ocl.c | 2 +- 6 files changed, 162 insertions(+), 37 deletions(-) diff --git a/kernel/alexkarnew.cl b/kernel/alexkarnew.cl index 43486c3e..488860ab 100644 --- a/kernel/alexkarnew.cl +++ b/kernel/alexkarnew.cl @@ -28,9 +28,34 @@ * online backup system. */ -/* Backwards compatibility, if NFACTOR not defined, default to 1024 scrypt */ +/* 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 1024 +#define NFACTOR 10 #endif __constant uint ES[2] = { 0x00FF00FF, 0xFF00FF00 }; @@ -766,7 +791,7 @@ 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; y0)); + 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; y0)); + 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; y0)); + 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<(NFACTOR/LOOKUP_GAP); ++y) + for(uint y=0; y<(N[NFACTOR]/LOOKUP_GAP); ++y) { for(uint z=0; zlookup_gap, (unsigned int)cgpu->thread_concurrency, (int)clState->wsize,(unsigned int)nfactor); + cgpu->lookup_gap, (unsigned int)cgpu->thread_concurrency, (int)clState->wsize, (unsigned int) opt_nfactor); applog(LOG_DEBUG, "Setting worksize to %d", (int)(clState->wsize)); if (clState->vwidth > 1) From 11cf733edee96556989e5a179aa3b04e2f38bafc Mon Sep 17 00:00:00 2001 From: Noel Maersk Date: Wed, 26 Feb 2014 04:32:49 +0200 Subject: [PATCH 11/21] misc: fix spaces after closing parenthesis. sed -i 's/) $/)/' *.cl --- kernel/alexkarnew.cl | 8 ++++---- kernel/alexkarold.cl | 8 ++++---- kernel/ckolivas.cl | 8 ++++---- kernel/psw.cl | 8 ++++---- kernel/zuikkis.cl | 6 +++--- 5 files changed, 19 insertions(+), 19 deletions(-) diff --git a/kernel/alexkarnew.cl b/kernel/alexkarnew.cl index 488860ab..ce9a629b 100644 --- a/kernel/alexkarnew.cl +++ b/kernel/alexkarnew.cl @@ -797,7 +797,7 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup) #pragma unroll for(uint z=0; z Date: Thu, 6 Mar 2014 20:55:56 +0200 Subject: [PATCH 12/21] algorithm: initial structure definition. --- Makefile.am | 1 + algorithm.c | 17 +++++++++++++++++ algorithm.h | 10 ++++++++++ 3 files changed, 28 insertions(+) create mode 100644 algorithm.c create mode 100644 algorithm.h diff --git a/Makefile.am b/Makefile.am index 1b753f36..b9804927 100644 --- a/Makefile.am +++ b/Makefile.am @@ -41,6 +41,7 @@ sgminer_SOURCES += driver-opencl.c driver-opencl.h sgminer_SOURCES += ocl.c ocl.h sgminer_SOURCES += findnonce.c findnonce.h sgminer_SOURCES += adl.c adl.h adl_functions.h +sgminer_SOURCES += algorithm.c algorithm.h sgminer_SOURCES += scrypt.c scrypt.h sgminer_SOURCES += kernel/*.cl diff --git a/algorithm.c b/algorithm.c new file mode 100644 index 00000000..f2c6a105 --- /dev/null +++ b/algorithm.c @@ -0,0 +1,17 @@ +/* + * 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 + +typedef struct algorithm_t { + char name[32]; /* Human-readable identifier */ + uint8_t nfactor; /* N factor (CPU/Memory tradeoff parameter) */ +} algorithm_t; diff --git a/algorithm.h b/algorithm.h new file mode 100644 index 00000000..10444e35 --- /dev/null +++ b/algorithm.h @@ -0,0 +1,10 @@ +#ifndef ALGORITHM_H +#define ALGORITHM_H + + +/* Describes the Scrypt parameters and hashing functions used to mine + * a specific coin. + */ +typedef struct algorithm_t algorithm_t; + +#endif /* ALGORITHM_H */ From 50a792f2cd8a8686f4083563b6477c364c9a79d5 Mon Sep 17 00:00:00 2001 From: Noel Maersk Date: Thu, 6 Mar 2014 21:22:33 +0200 Subject: [PATCH 13/21] algorithm: initial set_algorithm() and set_algorithm_nfactor(). --- algorithm.c | 18 +++++++++++++++++- algorithm.h | 7 +++++++ 2 files changed, 24 insertions(+), 1 deletion(-) diff --git a/algorithm.c b/algorithm.c index f2c6a105..f1f61a4b 100644 --- a/algorithm.c +++ b/algorithm.c @@ -10,8 +10,24 @@ #include "algorithm.h" #include +#include typedef struct algorithm_t { - char name[32]; /* Human-readable identifier */ + char name[20]; /* Human-readable identifier */ uint8_t nfactor; /* N factor (CPU/Memory tradeoff parameter) */ } algorithm_t; + +void set_algorithm(algorithm_t* algo, char* newname) { + strncpy(algo->name, newname, sizeof(algo->name)); + algo->name[sizeof(algo->name) - 1] = '\0'; + + if (strcmp(algo->name, "adaptive-nfactor") == 0) { + algo->nfactor = 11; + } else { + algo->nfactor = 10; + } +} + +void set_algorithm_nfactor(algorithm_t* algo, uint8_t nfactor) { + algo->nfactor = nfactor; +} diff --git a/algorithm.h b/algorithm.h index 10444e35..8d89eaef 100644 --- a/algorithm.h +++ b/algorithm.h @@ -1,10 +1,17 @@ #ifndef ALGORITHM_H #define ALGORITHM_H +#include /* Describes the Scrypt parameters and hashing functions used to mine * a specific coin. */ typedef struct algorithm_t algorithm_t; +/* Set default parameters based on name. */ +void set_algorithm(algorithm_t* algo, char* name); + +/* Set to specific N factor. */ +void set_algorithm_nfactor(algorithm_t* algo, uint8_t nfactor); + #endif /* ALGORITHM_H */ From 92b7770212389241568035452410825bc24114f9 Mon Sep 17 00:00:00 2001 From: Noel Maersk Date: Thu, 6 Mar 2014 21:59:51 +0200 Subject: [PATCH 14/21] config: add `--algorithm` option and documentation. Doc in `doc/configuration.md` (has to be started sometime, no?). Configuration function has to be lamely-named set_algo(), because set_algorithm() is already declared in algorithm.h (prevent namespace conflict). algorithm has to be added as global variable due to the way the callback is done (by CCAN/opt, which in itself is nice). This can be cleaned up significantly by (at least) introducing a global configuration struct, but there is no reason to do it now just for this - better a wholesale manana. --- algorithm.c | 4 ++-- algorithm.h | 4 ++-- doc/configuration.md | 26 ++++++++++++++++++++++++++ miner.h | 2 ++ sgminer.c | 18 +++++++++++++++++- 5 files changed, 49 insertions(+), 5 deletions(-) create mode 100644 doc/configuration.md diff --git a/algorithm.c b/algorithm.c index f1f61a4b..90f53ddf 100644 --- a/algorithm.c +++ b/algorithm.c @@ -17,7 +17,7 @@ typedef struct algorithm_t { uint8_t nfactor; /* N factor (CPU/Memory tradeoff parameter) */ } algorithm_t; -void set_algorithm(algorithm_t* algo, char* newname) { +void set_algorithm(algorithm_t* algo, const char* newname) { strncpy(algo->name, newname, sizeof(algo->name)); algo->name[sizeof(algo->name) - 1] = '\0'; @@ -28,6 +28,6 @@ void set_algorithm(algorithm_t* algo, char* newname) { } } -void set_algorithm_nfactor(algorithm_t* algo, uint8_t nfactor) { +void set_algorithm_nfactor(algorithm_t* algo, const uint8_t nfactor) { algo->nfactor = nfactor; } diff --git a/algorithm.h b/algorithm.h index 8d89eaef..b0749958 100644 --- a/algorithm.h +++ b/algorithm.h @@ -9,9 +9,9 @@ typedef struct algorithm_t algorithm_t; /* Set default parameters based on name. */ -void set_algorithm(algorithm_t* algo, char* name); +void set_algorithm(algorithm_t* algo, const char* name); /* Set to specific N factor. */ -void set_algorithm_nfactor(algorithm_t* algo, uint8_t nfactor); +void set_algorithm_nfactor(algorithm_t* algo, const uint8_t nfactor); #endif /* ALGORITHM_H */ diff --git a/doc/configuration.md b/doc/configuration.md new file mode 100644 index 00000000..8686d533 --- /dev/null +++ b/doc/configuration.md @@ -0,0 +1,26 @@ +# Configuration and command-line options + +*Work in progress!* + + +## Config-file and CLI options + +### algorithm + +Allows choosing between the few mining algorithms for incompatible +cryptocurrencies. + +Requires a string. + +Currently 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 N-factor scrypt parameter. + +Requires an unsigned integer. diff --git a/miner.h b/miner.h index 589bed5f..81dea7a9 100644 --- a/miner.h +++ b/miner.h @@ -1018,6 +1018,8 @@ extern bool fulltest(const unsigned char *hash, const unsigned char *target); extern int opt_queue; extern int opt_scantime; extern int opt_expiry; + +extern char* opt_algorithm; extern int opt_nfactor; extern cglock_t control_lock; diff --git a/sgminer.c b/sgminer.c index e3301337..6c5c265d 100644 --- a/sgminer.c +++ b/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" #if defined(unix) || defined(__APPLE__) @@ -92,7 +94,11 @@ int opt_log_interval = 5; int opt_queue = 1; int opt_scantime = 7; int opt_expiry = 28; + +char* opt_algorithm; +algorithm_t* algorithm; int opt_nfactor = 10; + static const bool opt_time = true; unsigned long long global_hashrate; unsigned long global_quota_gcd = 1; @@ -1005,6 +1011,13 @@ static void load_temp_cutoffs() } } +static char *set_algo(const char *arg) +{ + set_algorithm(algorithm, arg); + + return NULL; +} + static char *set_api_allow(const char *arg) { opt_set_charp(arg, &opt_api_allow); @@ -1054,6 +1067,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 to 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]"), @@ -1108,7 +1124,7 @@ static struct opt_table opt_config_table[] = { #endif OPT_WITH_ARG("--nfactor", set_int_0_to_9999, opt_show_intval, &opt_nfactor, - "Set scrypt nfactor, default: 10. Currently use 11 for vertcoin!"), + "Set scrypt N-factor parameter, default: 10. Currently use 11 for vertcoin!"), OPT_WITHOUT_ARG("--debug|-D", enable_debug, &opt_debug, "Enable debug output"), From a0c52bf67c378dff4ca1956a37e1c273740f1aed Mon Sep 17 00:00:00 2001 From: Noel Maersk Date: Thu, 6 Mar 2014 22:41:53 +0200 Subject: [PATCH 15/21] config: introduce set_nfactor() and use it to call set_algorithm_nfactor(). Had two bugs: 1. Will not compile due to unknown algorithm_t size. 2. nfactor is set to 0 (bad calling, fix later). So squashed two commits: 1. algorithm: move algorithm_t definition to header. 2. config: if --nfactor is specified, properly set algorithm->nfactor. --- algorithm.c | 5 ----- algorithm.h | 5 ++++- sgminer.c | 24 ++++++++++++++++-------- 3 files changed, 20 insertions(+), 14 deletions(-) diff --git a/algorithm.c b/algorithm.c index 90f53ddf..ba3d9cba 100644 --- a/algorithm.c +++ b/algorithm.c @@ -12,11 +12,6 @@ #include #include -typedef struct algorithm_t { - char name[20]; /* Human-readable identifier */ - uint8_t nfactor; /* N factor (CPU/Memory tradeoff parameter) */ -} algorithm_t; - void set_algorithm(algorithm_t* algo, const char* newname) { strncpy(algo->name, newname, sizeof(algo->name)); algo->name[sizeof(algo->name) - 1] = '\0'; diff --git a/algorithm.h b/algorithm.h index b0749958..0c73d9bb 100644 --- a/algorithm.h +++ b/algorithm.h @@ -6,7 +6,10 @@ /* Describes the Scrypt parameters and hashing functions used to mine * a specific coin. */ -typedef struct algorithm_t algorithm_t; +typedef struct _algorithm_t { + char name[20]; /* Human-readable identifier */ + uint8_t nfactor; /* N factor (CPU/Memory tradeoff parameter) */ +} algorithm_t; /* Set default parameters based on name. */ void set_algorithm(algorithm_t* algo, const char* name); diff --git a/sgminer.c b/sgminer.c index 6c5c265d..aaeb3828 100644 --- a/sgminer.c +++ b/sgminer.c @@ -95,8 +95,8 @@ int opt_queue = 1; int opt_scantime = 7; int opt_expiry = 28; -char* opt_algorithm; -algorithm_t* algorithm; +char *opt_algorithm; +algorithm_t *algorithm; int opt_nfactor = 10; static const bool opt_time = true; @@ -1018,6 +1018,12 @@ static char *set_algo(const char *arg) return NULL; } +static char *set_nfactor(const char *arg) +{ + set_algorithm_nfactor(algorithm, (uint8_t)atoi(arg)); + return NULL; +} + static char *set_api_allow(const char *arg) { opt_set_charp(arg, &opt_api_allow); @@ -1069,7 +1075,7 @@ static char *set_null(const char __maybe_unused *arg) static struct opt_table opt_config_table[] = { OPT_WITH_ARG("--algorithm", set_algo, NULL, NULL, - "Set mining algorithm to most common defaults, default: static"), + "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]"), @@ -1122,9 +1128,6 @@ static struct opt_table opt_config_table[] = { opt_set_bool, &opt_compact, "Use compact display without per device statistics"), #endif - OPT_WITH_ARG("--nfactor", - set_int_0_to_9999, opt_show_intval, &opt_nfactor, - "Set scrypt N-factor parameter, default: 10. Currently use 11 for vertcoin!"), OPT_WITHOUT_ARG("--debug|-D", enable_debug, &opt_debug, "Enable debug output"), @@ -1228,13 +1231,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, @@ -7764,7 +7770,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"); @@ -7842,6 +7848,8 @@ int main(int argc, char *argv[]) strcat(sgminer_path, "\\"); #endif + algorithm = (algorithm_t *)alloca(sizeof(algorithm_t)); + devcursor = 8; logstart = devcursor + 1; logcursor = logstart + 1; From c6a27709f8d2cc47af656baaa8b809cb54b0f16f Mon Sep 17 00:00:00 2001 From: Noel Maersk Date: Fri, 7 Mar 2014 01:29:55 +0200 Subject: [PATCH 16/21] core: use global algorithm structure instead of opt_nfactor. Also squashed: config: add log messages to set_algo() and set_nfactor(). algorithm: use set_algorithm_nfactor() when setting default nfactor in set_algorithm(). Otherwise algorithm->n defaults to 0. P.S. Did I already mention how this could have been C++?.. --- algorithm.c | 9 +++++++-- algorithm.h | 5 +++-- miner.h | 6 ++++-- ocl.c | 18 +++++++++++------- scrypt.c | 24 +++++++++++++----------- sgminer.c | 5 ++++- 6 files changed, 42 insertions(+), 25 deletions(-) diff --git a/algorithm.c b/algorithm.c index ba3d9cba..7399fe02 100644 --- a/algorithm.c +++ b/algorithm.c @@ -17,12 +17,17 @@ void set_algorithm(algorithm_t* algo, const char* newname) { algo->name[sizeof(algo->name) - 1] = '\0'; if (strcmp(algo->name, "adaptive-nfactor") == 0) { - algo->nfactor = 11; + set_algorithm_nfactor(algo, 11); } else { - algo->nfactor = 10; + 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; } diff --git a/algorithm.h b/algorithm.h index 0c73d9bb..b6053994 100644 --- a/algorithm.h +++ b/algorithm.h @@ -7,8 +7,9 @@ * a specific coin. */ typedef struct _algorithm_t { - char name[20]; /* Human-readable identifier */ - uint8_t nfactor; /* N factor (CPU/Memory tradeoff parameter) */ + 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. */ diff --git a/miner.h b/miner.h index 81dea7a9..c932ab1b 100644 --- a/miner.h +++ b/miner.h @@ -3,6 +3,8 @@ #include "config.h" +#include "algorithm.h" + #include #include #include @@ -1019,8 +1021,8 @@ extern int opt_queue; extern int opt_scantime; extern int opt_expiry; -extern char* opt_algorithm; -extern int opt_nfactor; +extern char *opt_algorithm; +extern algorithm_t *algorithm; extern cglock_t control_lock; extern pthread_mutex_t hash_lock; diff --git a/ocl.c b/ocl.c index a456a6a7..68e6ed8b 100644 --- a/ocl.c +++ b/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) @@ -226,9 +232,6 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize) cl_uint numDevices; cl_int status; - /* Scrypt CPU/Memory cost parameter */ - cl_uint N = (1 << opt_nfactor); - status = clGetPlatformIDs(0, NULL, &numPlatforms); if (status != CL_SUCCESS) { applog(LOG_ERR, "Error %d: Getting Platforms. (clGetPlatformsIDs)", status); @@ -485,7 +488,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize) if (!cgpu->opt_tc) { unsigned int sixtyfours; - sixtyfours = cgpu->max_alloc / 131072 / 64 / (N/1024) - 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; @@ -525,7 +528,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize) if (clState->goffset) strcat(binaryfilename, "g"); - sprintf(numbuf, "lg%utc%unf%u", cgpu->lookup_gap, (unsigned int)cgpu->thread_concurrency,opt_nfactor); + 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); @@ -592,7 +595,7 @@ build: char *CompilerOptions = (char *)calloc(1, 256); 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, (unsigned int) opt_nfactor); + 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) @@ -781,7 +784,8 @@ built: return NULL; } - size_t ipt = (N / cgpu->lookup_gap + (N % 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 diff --git a/scrypt.c b/scrypt.c index 3b2e74de..3a48ea14 100644 --- a/scrypt.c +++ b/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 - scratchpad size needs to be at least 63 + (128 * r * p) + (256 * r + 64) + (128 * r * N) bytes +/* cpu and memory intensive function to transform a 80 byte buffer into + * 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_n_1_1_256_sp(const uint32_t* input, char* scratchpad, uint32_t *ostate, const cl_uint n) +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_n_1_1_256_sp(const uint32_t* input, char* scratchpad, uint32_ PBKDF2_SHA256_80_128(input, X); - for (i = 0; i < n; 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_n_1_1_256_sp(const uint32_t* input, char* scratchpad, uint32_ salsa20_8(&X[0], &X[16]); salsa20_8(&X[16], &X[0]); } - for (i = 0; i < n; i += 2) { - j = X[16] & (n-1); + for (i = 0; i < algorithm->n; i += 2) { + 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_n_1_1_256_sp(const uint32_t* input, char* scratchpad, uint32_ salsa20_8(&X[0], &X[16]); salsa20_8(&X[16], &X[0]); - j = X[16] & (n-1); + j = X[16] & (algorithm->n-1); p2 = (uint64_t *)(&V[j * 32]); for(k = 0; k < 16; k++) p1[k] ^= p2[k]; @@ -412,8 +414,8 @@ void scrypt_regenhash(struct work *work) be32enc_vect(data, (const uint32_t *)work->data, 19); data[19] = htobe32(*nonce); - scratchbuf = (char *)alloca((1 << opt_nfactor) * 128 + 512); - scrypt_n_1_1_256_sp(data, scratchbuf, ohash, (1 << opt_nfactor)); + scratchbuf = (char *)alloca(algorithm->n * 128 + 512); + scrypt_n_1_1_256_sp(data, scratchbuf, ohash); flip32(ohash, ohash); } @@ -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_n_1_1_256_sp(data, scratchbuf, ohash, (1 << opt_nfactor)); + 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_n_1_1_256_sp(data, scratchbuf, ostate, (1 << opt_nfactor)); + scrypt_n_1_1_256_sp(data, scratchbuf, ostate, algorithm->n); tmp_hash7 = be32toh(ostate[7]); if (unlikely(tmp_hash7 <= Htarg)) { diff --git a/sgminer.c b/sgminer.c index aaeb3828..d759c45d 100644 --- a/sgminer.c +++ b/sgminer.c @@ -97,7 +97,6 @@ int opt_expiry = 28; char *opt_algorithm; algorithm_t *algorithm; -int opt_nfactor = 10; static const bool opt_time = true; unsigned long long global_hashrate; @@ -1014,6 +1013,7 @@ 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; } @@ -1021,6 +1021,9 @@ static char *set_algo(const char *arg) 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; } From a3f9b24c697d8b003ea6772b0cf48c88a88f23fd Mon Sep 17 00:00:00 2001 From: Noel Maersk Date: Fri, 7 Mar 2014 02:21:09 +0200 Subject: [PATCH 17/21] core: set default algorithm for when neither --algorithm nor --nfactor are provided. --- sgminer.c | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sgminer.c b/sgminer.c index d759c45d..f4a1790e 100644 --- a/sgminer.c +++ b/sgminer.c @@ -7851,7 +7851,9 @@ 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, "default"); devcursor = 8; logstart = devcursor + 1; From 4ed13e44a3ecb52f5740443df340397812fb82c7 Mon Sep 17 00:00:00 2001 From: Noel Maersk Date: Fri, 7 Mar 2014 02:21:42 +0200 Subject: [PATCH 18/21] doc: update configuration.md with latest on --algorithm and --nfactor. Should close https://github.com/veox/sgminer/issues/126 --- doc/configuration.md | 18 ++++++++++++++---- 1 file changed, 14 insertions(+), 4 deletions(-) diff --git a/doc/configuration.md b/doc/configuration.md index 8686d533..9d99a718 100644 --- a/doc/configuration.md +++ b/doc/configuration.md @@ -10,9 +10,11 @@ Allows choosing between the few mining algorithms for incompatible cryptocurrencies. -Requires a string. +*Argument:* string -Currently supported: +*Default:* `default` + +*Supported:* * `adaptive-nfactor` - Vertcoin-style adaptive N-factor scrypt. N-factor defaults to 11. @@ -21,6 +23,14 @@ N-factor defaults to 11. ### nfactor -Overrides the default N-factor scrypt parameter. +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 -Requires an unsigned integer. +*TODO* From 201036dd923e52e4e7ce46a540485f50b6ae320c Mon Sep 17 00:00:00 2001 From: unknown Date: Sun, 9 Mar 2014 11:29:53 +0100 Subject: [PATCH 19/21] Added algorithm.* and inttypes.h for MSVS build --- winbuild/dist/include/inttypes.h | 305 +++++++++++++++++++++++++++++++ winbuild/sgminer.vcxproj | 2 + winbuild/sgminer.vcxproj.filters | 6 + 3 files changed, 313 insertions(+) create mode 100644 winbuild/dist/include/inttypes.h diff --git a/winbuild/dist/include/inttypes.h b/winbuild/dist/include/inttypes.h new file mode 100644 index 00000000..4b3828a2 --- /dev/null +++ b/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_ ] diff --git a/winbuild/sgminer.vcxproj b/winbuild/sgminer.vcxproj index 7bcb78db..fb9a10d5 100644 --- a/winbuild/sgminer.vcxproj +++ b/winbuild/sgminer.vcxproj @@ -241,6 +241,7 @@ exit 0 true + @@ -268,6 +269,7 @@ exit 0 + diff --git a/winbuild/sgminer.vcxproj.filters b/winbuild/sgminer.vcxproj.filters index 0ed5cb60..a2fd51fc 100644 --- a/winbuild/sgminer.vcxproj.filters +++ b/winbuild/sgminer.vcxproj.filters @@ -92,6 +92,9 @@ Libraries\jansson + + Source Files + @@ -163,6 +166,9 @@ Header Files + + Header Files + From 0e37fc24b0908c72d6b6c502cc6bc636fb12d6cf Mon Sep 17 00:00:00 2001 From: unknown Date: Thu, 13 Mar 2014 23:02:51 +0100 Subject: [PATCH 20/21] Added algorrithm name to "coin" API command. Changed default algo name to "scrypt". --- algorithm.c | 4 ++-- api.c | 2 +- sgminer.c | 2 +- 3 files changed, 4 insertions(+), 4 deletions(-) diff --git a/algorithm.c b/algorithm.c index 7399fe02..bafb1dcb 100644 --- a/algorithm.c +++ b/algorithm.c @@ -17,9 +17,9 @@ void set_algorithm(algorithm_t* algo, const char* newname) { algo->name[sizeof(algo->name) - 1] = '\0'; if (strcmp(algo->name, "adaptive-nfactor") == 0) { - set_algorithm_nfactor(algo, 11); + set_algorithm_nfactor(algo, 11); } else { - set_algorithm_nfactor(algo, 10); + set_algorithm_nfactor(algo, 10); } return; diff --git a/api.c b/api.c index 3b0cff47..e0274726 100644 --- a/api.c +++ b/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); diff --git a/sgminer.c b/sgminer.c index f4a1790e..81b4cd5a 100644 --- a/sgminer.c +++ b/sgminer.c @@ -7853,7 +7853,7 @@ int main(int argc, char *argv[]) /* Default algorithm specified in algorithm.c ATM */ algorithm = (algorithm_t *)alloca(sizeof(algorithm_t)); - set_algorithm(algorithm, "default"); + set_algorithm(algorithm, "scrypt"); devcursor = 8; logstart = devcursor + 1; From 0db99d66c00c8225703703a7360d5c1497404bba Mon Sep 17 00:00:00 2001 From: Noel Maersk Date: Fri, 14 Mar 2014 17:05:23 +0200 Subject: [PATCH 21/21] algorithm: add nfactor aliases "adaptive-n-factor" and "nscrypt". Also change whitespace back to "4 spaces". --- algorithm.c | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/algorithm.c b/algorithm.c index bafb1dcb..cf09f60c 100644 --- a/algorithm.c +++ b/algorithm.c @@ -16,10 +16,12 @@ 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-nfactor") == 0) { - set_algorithm_nfactor(algo, 11); + 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); + set_algorithm_nfactor(algo, 10); } return;