v3.12.5

README.md

@@ -12,14 +12,23 @@ a false positive, they are flagged simply because they are cryptocurrency
 miners. The source code is open for anyone to inspect. If you don't trust 
 the software, don't use it.
 
+
+New thread:
+
+https://bitcointalk.org/index.php?topic=5226770.msg53865575#msg53865575
+
+Old thread:
+
 https://bitcointalk.org/index.php?topic=1326803.0
 
 mailto://jayddee246@gmail.com
 
+This note is to confirm that bitcointalk users JayDDee and joblo are the
+same person.
 
 I created a new BCT user JayDDee to match my github user id.
-Over timme all discussion will be migrated using the new user.
-This note is to confirm JayDDee is the sae person as joblo.
+The old thread has been locked but still contains useful information for
+reading.
 
 See file RELEASE_NOTES for change log and INSTALL_LINUX or INSTALL_WINDOWS
 for compile instructions.
@@ -143,6 +152,27 @@ Supported Algorithms
                           yespower-b2b  generic yespower + blake2b
                           zr5           Ziftr
 
+Many variations of scrypt based algos can be mine by specifying their
+parameters:
+
+scryptn2: --algo scrypt --param-n 1048576
+
+cpupower: --algo yespower --param-key "CPUpower: The number of CPU working or available for proof-of-work mining"
+
+power2b: --algo yespower-b2b --param-n 2048 --param-r 32 --param-key "Now I am become Death, the destroyer of worlds"
+
+sugarchain: --algo yespower --param-n 2048 -param-r 32 --param-key "Satoshi Nakamoto 31/Oct/2008 Proof-of-work is essentially one-CPU-one-vote"
+
+yespoweriots: --a yespower --param-n 2048 --param-key "Iots is committed to the development of IOT"
+
+yespowerlitb: --algo yespower --param-n 2048 --param-r 32 --param-key "LITBpower: The number of LITB working or available for proof-of-work mini"
+
+yespoweric: --algo yespower --param-n 2048 --param-r 32 --param-key "IsotopeC" 
+
+yespowerurx: --algo yespower --param-n 2048 --param-r 32 --param-key "UraniumX"
+
+yespowerltncg: --algo yespower --param-n 2048 --param-r 32 --param-key "LTNCGYES"
+
 Errata
 ------
 

RELEASE_NOTES

@@ -65,9 +65,79 @@ If not what makes it happen or not happen?
 Change Log
 ----------
 
+v3.12.5
+
+Issues #246 & #251: fixed incorrect share diff for stratum and getwork,
+fixed incorrect target diff for getwork. Stats should now be correct for
+getwork as well as stratum.
+
+Getwork: reduce stale blocks, faster response to new work.
+
+Added ntime to new job/work logs.
+
+README.md now lists the parameters for yespower variations that don't have
+a specific algo name.
+
+v3.12.4.6
+
+Issue #246: fixed getwork repeated new block logs with same height. New work
+for the same block is now reported as "New work" instead of New block".
+Also added a check that work is new before generating "New work" log.
+
+Added target diff to getwork new block log.
+
+Changed share ratio in share result log to simple fraction, no longer %.
+
+Added debug log to display mininginfo, use -D.
+
+v3.12.4.5
+
+Issue #246: better stale share detection for getwork, and enhanced logging
+of stale shares for stratum & getwork.
+
+Issue #251: fixed incorrect share difficulty and share ratio in share
+result log.
+
+Changed submit log to include share diff and block height.
+
+Small cosmetic changes to logs. 
+
+v3.12.4.4
+
+Issue #246: Fixed net hashrate in getwork block log,
+            removed duplicate getwork block log, 
+            other small tweaks to stats logs for getwork.
+
+Issue #248: Fixed chronic stale shares with scrypt:1048576 (scryptn2). 
+
+v3.12.4.3
+
+Fixed segfault in new block log for getwork.
+
+Disabled silent discarding of stale work after the submit is logged.
+
+v3.12.4.2
+
+Issue #245: fixed getwork stale shares, solo mining with getwork now works.
+
+Issue #246: implemented block and summary logs for getwork.
+
+v3.12.4.1
+
+Issue #245: fix scantime when mining solo with getwork.
+
+Added debug logs for creation of stratum and longpoll threads, use -D to
+enable.
+
+v3.12.4
+
+Issue #244: Change longpoll to ignore job id.
+
+Lyra2rev2 AVX2 +3%, AVX512 +6%.
+
 v3.12.3.1
 
-Issue #241: Fixed regression that broke GBT in v3.11.2.
+Issue #241: Fixed regression that broke coinbase address in v3.11.7.
 
 v3.12.3
 

algo/lyra2/lyra2-gate.c

@@ -94,12 +94,12 @@ bool lyra2rev2_thread_init()
    const int64_t ROW_LEN_BYTES = ROW_LEN_INT64 * 8;
 
    int size = (int64_t)ROW_LEN_BYTES * 4; // nRows;
-#if defined (LYRA2REV2_8WAY)
+#if defined (LYRA2REV2_16WAY)
    l2v2_wholeMatrix = _mm_malloc( 2 * size, 64 );   // 2 way
-   init_lyra2rev2_8way_ctx();;
-#elif defined (LYRA2REV2_4WAY)
+   init_lyra2rev2_16way_ctx();;
+#elif defined (LYRA2REV2_8WAY)
    l2v2_wholeMatrix = _mm_malloc( size, 64 );
-   init_lyra2rev2_4way_ctx();;
+   init_lyra2rev2_8way_ctx();;
 #else
    l2v2_wholeMatrix = _mm_malloc( size, 64 );
    init_lyra2rev2_ctx();
@@ -109,12 +109,12 @@ bool lyra2rev2_thread_init()
 
 bool register_lyra2rev2_algo( algo_gate_t* gate )
 {
-#if defined (LYRA2REV2_8WAY)
+#if defined (LYRA2REV2_16WAY)
+  gate->scanhash  = (void*)&scanhash_lyra2rev2_16way;
+  gate->hash      = (void*)&lyra2rev2_16way_hash;
+#elif defined (LYRA2REV2_8WAY)
   gate->scanhash  = (void*)&scanhash_lyra2rev2_8way;
   gate->hash      = (void*)&lyra2rev2_8way_hash;
-#elif defined (LYRA2REV2_4WAY)
-  gate->scanhash  = (void*)&scanhash_lyra2rev2_4way;
-  gate->hash      = (void*)&lyra2rev2_4way_hash;
 #else
   gate->scanhash  = (void*)&scanhash_lyra2rev2;
   gate->hash      = (void*)&lyra2rev2_hash;

algo/lyra2/lyra2-gate.h

@@ -51,30 +51,32 @@ bool init_lyra2rev3_ctx();
 //////////////////////////////////
 
 #if defined(__AVX512F__) && defined(__AVX512VL__) && defined(__AVX512DQ__) && defined(__AVX512BW__)
-  #define LYRA2REV2_8WAY 1
+  #define LYRA2REV2_16WAY 1
 #elif defined(__AVX2__)
-  #define LYRA2REV2_4WAY 1
+  #define LYRA2REV2_8WAY 1
 #endif
 
 extern __thread uint64_t* l2v2_wholeMatrix;
 
 bool register_lyra2rev2_algo( algo_gate_t* gate );
 
-#if defined(LYRA2REV2_8WAY)
+#if defined(LYRA2REV2_16WAY)
+
+void lyra2rev2_16way_hash( void *state, const void *input );
+int scanhash_lyra2rev2_16way( struct work *work, uint32_t max_nonce,
+                             uint64_t *hashes_done, struct thr_info *mythr );
+bool init_lyra2rev2_16way_ctx();
+
+#elif defined(LYRA2REV2_8WAY)
 
 void lyra2rev2_8way_hash( void *state, const void *input );
 int scanhash_lyra2rev2_8way( struct work *work, uint32_t max_nonce,
                              uint64_t *hashes_done, struct thr_info *mythr );
 bool init_lyra2rev2_8way_ctx();
 
-#elif defined(LYRA2REV2_4WAY)
-
-void lyra2rev2_4way_hash( void *state, const void *input );
-int scanhash_lyra2rev2_4way( struct work *work, uint32_t max_nonce,
-                             uint64_t *hashes_done, struct thr_info *mythr );
-bool init_lyra2rev2_4way_ctx();
 
 #else
+
 void lyra2rev2_hash( void *state, const void *input );
 int scanhash_lyra2rev2( struct work *work, uint32_t max_nonce,
                         uint64_t *hashes_done, struct thr_info *mythr );

algo/lyra2/lyra2rev2-4way.c

@@ -8,12 +8,30 @@
 #include "algo/cubehash/cube-hash-2way.h"
 
 
-#if 0
-void lyra2rev2_8way_hash( void *state, const void *input )
+#if defined (LYRA2REV2_16WAY)
+
+typedef struct {
+   blake256_16way_context    blake;
+   keccak256_8way_context    keccak;
+   cubehashParam             cube;
+   skein256_8way_context     skein;
+   bmw256_16way_context      bmw;
+} lyra2v2_16way_ctx_holder __attribute__ ((aligned (64)));
+
+static lyra2v2_16way_ctx_holder l2v2_16way_ctx;
+
+bool init_lyra2rev2_16way_ctx()
 {
-   uint32_t vhash[8*8] __attribute__ ((aligned (128)));
-   uint32_t vhashA[8*8] __attribute__ ((aligned (64)));
-   uint32_t vhashB[8*8] __attribute__ ((aligned (64)));
+   keccak256_8way_init( &l2v2_16way_ctx.keccak );
+   cubehashInit( &l2v2_16way_ctx.cube, 256, 16, 32 );
+   skein256_8way_init( &l2v2_16way_ctx.skein );
+   bmw256_16way_init( &l2v2_16way_ctx.bmw );
+   return true;
+}
+
+void lyra2rev2_16way_hash( void *state, const void *input )
+{
+   uint32_t vhash[8*16] __attribute__ ((aligned (128)));
    uint32_t hash0[8] __attribute__ ((aligned (64)));
    uint32_t hash1[8] __attribute__ ((aligned (64)));
    uint32_t hash2[8] __attribute__ ((aligned (64)));
@@ -22,19 +40,42 @@ void lyra2rev2_8way_hash( void *state, const void *input )
    uint32_t hash5[8] __attribute__ ((aligned (64)));
    uint32_t hash6[8] __attribute__ ((aligned (64)));
    uint32_t hash7[8] __attribute__ ((aligned (64)));
-   lyra2v2_8way_ctx_holder ctx __attribute__ ((aligned (64)));
-   memcpy( &ctx, &l2v2_8way_ctx, sizeof(l2v2_8way_ctx) );
+   uint32_t hash8[8] __attribute__ ((aligned (64)));
+   uint32_t hash9[8] __attribute__ ((aligned (64)));
+   uint32_t hash10[8] __attribute__ ((aligned (64)));
+   uint32_t hash11[8] __attribute__ ((aligned (64)));
+   uint32_t hash12[8] __attribute__ ((aligned (64)));
+   uint32_t hash13[8] __attribute__ ((aligned (64)));
+   uint32_t hash14[8] __attribute__ ((aligned (64)));
+   uint32_t hash15[8] __attribute__ ((aligned (64)));
+   lyra2v2_16way_ctx_holder ctx __attribute__ ((aligned (64)));
+   memcpy( &ctx, &l2v2_16way_ctx, sizeof(l2v2_16way_ctx) );
 
-   blake256_8way_update( &ctx.blake, input + (64<<3), 16 );
-   blake256_8way_close( &ctx.blake, vhash );
+   blake256_16way_update( &ctx.blake, input + (64<<4), 16 );
+   blake256_16way_close( &ctx.blake, vhash );
 
-   rintrlv_8x32_8x64( vhashA, vhash, 256 );
+   dintrlv_16x32( hash0,  hash1,  hash2,  hash3,
+                  hash4,  hash5,  hash6,  hash7,
+                  hash8,  hash9,  hash10, hash11,
+                  hash12, hash13, hash14, hash15, vhash, 256 );
 
-   keccak256_8way_update( &ctx.keccak, vhashA, 32 );
+   intrlv_8x64( vhash, hash0, hash1, hash2, hash3,
+                       hash4, hash5, hash6, hash7, 256 );
+
+   keccak256_8way_update( &ctx.keccak, vhash, 32 );
    keccak256_8way_close( &ctx.keccak, vhash );
 
    dintrlv_8x64( hash0, hash1, hash2, hash3,
                  hash4, hash5, hash6, hash7, vhash, 256 );
+   intrlv_8x64( vhash, hash8,  hash9,  hash10, hash11,
+                       hash12, hash13, hash14, hash15, 256 );
+
+   keccak256_8way_init( &ctx.keccak );
+   keccak256_8way_update( &ctx.keccak, vhash, 32 );
+   keccak256_8way_close( &ctx.keccak, vhash );
+
+   dintrlv_8x64( hash8,  hash9,  hash10,  hash11,
+                 hash12, hash13, hash14, hash5, vhash, 256 );
 
    cubehash_full( &ctx.cube, (byte*) hash0,  256, (const byte*) hash0,  32 );
    cubehash_full( &ctx.cube, (byte*) hash1,  256, (const byte*) hash1,  32 );
@@ -44,13 +85,15 @@ void lyra2rev2_8way_hash( void *state, const void *input )
    cubehash_full( &ctx.cube, (byte*) hash5,  256, (const byte*) hash5,  32 );
    cubehash_full( &ctx.cube, (byte*) hash6,  256, (const byte*) hash6,  32 );
    cubehash_full( &ctx.cube, (byte*) hash7,  256, (const byte*) hash7,  32 );
+   cubehash_full( &ctx.cube, (byte*) hash8,  256, (const byte*) hash8,  32 );
+   cubehash_full( &ctx.cube, (byte*) hash9,  256, (const byte*) hash9,  32 );
+   cubehash_full( &ctx.cube, (byte*) hash10, 256, (const byte*) hash10, 32 );
+   cubehash_full( &ctx.cube, (byte*) hash11, 256, (const byte*) hash11, 32 );
+   cubehash_full( &ctx.cube, (byte*) hash12, 256, (const byte*) hash12, 32 );
+   cubehash_full( &ctx.cube, (byte*) hash13, 256, (const byte*) hash13, 32 );
+   cubehash_full( &ctx.cube, (byte*) hash14, 256, (const byte*) hash14, 32 );
+   cubehash_full( &ctx.cube, (byte*) hash15, 256, (const byte*) hash15, 32 );
 
-//   cube_4way_update_close( &ctx.cube, vhashA, vhashA, 32 );
-//   cube_4way_init( &ctx.cube, 256, 16, 32 );
-//   cube_4way_update_close( &ctx.cube, vhashB, vhashB, 32 );
-//
-//   dintrlv_4x128( hash0, hash1, hash2, hash3, vhashA, 256 );
-//   dintrlv_4x128( hash4, hash5, hash6, hash7, vhashB, 256 );
 
    intrlv_2x256( vhash, hash0, hash1, 256 );
    LYRA2REV2_2WAY( l2v2_wholeMatrix, vhash, 32, vhash, 32, 1, 4, 4 );
@@ -64,15 +107,36 @@ void lyra2rev2_8way_hash( void *state, const void *input )
    intrlv_2x256( vhash, hash6, hash7, 256 );
    LYRA2REV2_2WAY( l2v2_wholeMatrix, vhash, 32, vhash, 32, 1, 4, 4 );
    dintrlv_2x256( hash6, hash7, vhash, 256 );
+   intrlv_2x256( vhash, hash8, hash9, 256 );
+   LYRA2REV2_2WAY( l2v2_wholeMatrix, vhash, 32, vhash, 32, 1, 4, 4 );
+   dintrlv_2x256( hash8, hash9, vhash, 256 );
+   intrlv_2x256( vhash, hash10, hash11, 256 );
+   LYRA2REV2_2WAY( l2v2_wholeMatrix, vhash, 32, vhash, 32, 1, 4, 4 );
+   dintrlv_2x256( hash10, hash11, vhash, 256 );
+   intrlv_2x256( vhash, hash12, hash13, 256 );
+   LYRA2REV2_2WAY( l2v2_wholeMatrix, vhash, 32, vhash, 32, 1, 4, 4 );
+   dintrlv_2x256( hash12, hash13, vhash, 256 );
+   intrlv_2x256( vhash, hash14, hash15, 256 );
+   LYRA2REV2_2WAY( l2v2_wholeMatrix, vhash, 32, vhash, 32, 1, 4, 4 );
+   dintrlv_2x256( hash14, hash15, vhash, 256 );
 
-   intrlv_8x64( vhash, hash0, hash1, hash2, hash3, hash4, hash5, hash6,
-                hash7, 256 );
-
+   intrlv_8x64( vhash, hash0, hash1, hash2, hash3,
+                       hash4, hash5, hash6, hash7, 256 );
    skein256_8way_update( &ctx.skein, vhash, 32 );
    skein256_8way_close( &ctx.skein, vhash );
 
    dintrlv_8x64( hash0, hash1, hash2, hash3,
                  hash4, hash5, hash6, hash7, vhash, 256 );
+   intrlv_8x64( vhash, hash8,  hash9,  hash10, hash11, hash12,
+                       hash13, hash14, hash15, 256 );
+
+   skein256_8way_init( &ctx.skein );
+   skein256_8way_update( &ctx.skein, vhash, 32 );
+   skein256_8way_close( &ctx.skein, vhash );
+
+   dintrlv_8x64( hash8,  hash9,  hash10, hash11,
+                 hash12, hash13, hash14, hash15, vhash, 256 );
+
    
    cubehash_full( &ctx.cube, (byte*) hash0,  256, (const byte*) hash0, 32 );
    cubehash_full( &ctx.cube, (byte*) hash1,  256, (const byte*) hash1, 32 );
@@ -82,33 +146,78 @@ void lyra2rev2_8way_hash( void *state, const void *input )
    cubehash_full( &ctx.cube, (byte*) hash5,  256, (const byte*) hash5, 32 );
    cubehash_full( &ctx.cube, (byte*) hash6,  256, (const byte*) hash6, 32 );
    cubehash_full( &ctx.cube, (byte*) hash7,  256, (const byte*) hash7, 32 );
+   cubehash_full( &ctx.cube, (byte*) hash8,  256, (const byte*) hash8,  32 );
+   cubehash_full( &ctx.cube, (byte*) hash9,  256, (const byte*) hash9,  32 );
+   cubehash_full( &ctx.cube, (byte*) hash10, 256, (const byte*) hash10, 32 );
+   cubehash_full( &ctx.cube, (byte*) hash11, 256, (const byte*) hash11, 32 );
+   cubehash_full( &ctx.cube, (byte*) hash12, 256, (const byte*) hash12, 32 );
+   cubehash_full( &ctx.cube, (byte*) hash13, 256, (const byte*) hash13, 32 );
+   cubehash_full( &ctx.cube, (byte*) hash14, 256, (const byte*) hash14, 32 );
+   cubehash_full( &ctx.cube, (byte*) hash15, 256, (const byte*) hash15, 32 );
 
-//   cube_4way_init( &ctx.cube, 256, 16, 32 );
-//   cube_4way_update_close( &ctx.cube, vhashA, vhashA, 32 );
-//   cube_4way_init( &ctx.cube, 256, 16, 32 );
-//   cube_4way_update_close( &ctx.cube, vhashB, vhashB, 32 );
-//   
-//   dintrlv_4x128( hash0, hash1, hash2, hash3, vhashA, 256 );
-//   dintrlv_4x128( hash4, hash5, hash6, hash7, vhashB, 256 );
+   intrlv_16x32( vhash, hash0,  hash1,  hash2,  hash3,
+                        hash4,  hash5,  hash6,  hash7,
+                        hash8,  hash9,  hash10, hash11,
+                        hash12, hash13, hash14, hash15, 256 );
 
-   intrlv_8x32( vhash, hash0, hash1, hash2, hash3, hash4, hash5, hash6,
-                hash7, 256 );
-
-   bmw256_8way_update( &ctx.bmw, vhash, 32 );
-   bmw256_8way_close( &ctx.bmw, state );
+   bmw256_16way_update( &ctx.bmw, vhash, 32 );
+   bmw256_16way_close( &ctx.bmw, state );
 }
-#endif
 
+int scanhash_lyra2rev2_16way( struct work *work, const uint32_t max_nonce,
+                             uint64_t *hashes_done, struct thr_info *mythr )
+{
+   uint32_t hash[8*16] __attribute__ ((aligned (128)));
+   uint32_t vdata[20*16] __attribute__ ((aligned (64)));
+   uint32_t *hashd7 = &hash[7*16];
+   uint32_t lane_hash[8] __attribute__ ((aligned (32)));
+   uint32_t *pdata = work->data;
+   uint32_t *ptarget = work->target;
+   const uint32_t first_nonce = pdata[19];
+   const uint32_t last_nonce = max_nonce - 16;
+   uint32_t n = first_nonce;
+   const uint32_t targ32 = ptarget[7];
+   __m512i  *noncev = (__m512i*)vdata + 19;
+   const int thr_id = mythr->id;
+   const bool bench = opt_benchmark;
 
+   if ( bench )  ptarget[7] = 0x0000ff;
 
+   mm512_bswap32_intrlv80_16x32( vdata, pdata );
+   *noncev = _mm512_set_epi32( n+15, n+14, n+13, n+12, n+11, n+10, n+ 9, n+ 8,
+                               n+ 7, n+ 6, n+ 5, n+ 4, n+ 3, n+ 2, n+ 1, n );
+   blake256_16way_init( &l2v2_16way_ctx.blake );
+   blake256_16way_update( &l2v2_16way_ctx.blake, vdata, 64 );
 
-#if defined (LYRA2REV2_8WAY)
+   do
+   {
+      lyra2rev2_16way_hash( hash, vdata );
+
+      for ( int lane = 0; lane < 16; lane++ )
+      if ( unlikely( hashd7[lane] <= targ32 ) )
+      {
+         extr_lane_16x32( lane_hash, hash, lane, 256 );
+         if ( likely( valid_hash( lane_hash, ptarget ) && !bench ) )
+         {
+             pdata[19] = bswap_32( n + lane );
+             submit_lane_solution( work, lane_hash, mythr, lane );
+         }
+      }
+      *noncev = _mm512_add_epi32( *noncev, m512_const1_32( 16 ) );
+      n += 16;
+   } while ( likely( (n < last_nonce) && !work_restart[thr_id].restart ) );
+   pdata[19] = n;
+   *hashes_done = n - first_nonce;
+   return 0;
+}
+
+#elif defined (LYRA2REV2_8WAY)
 
 typedef struct {
    blake256_8way_context     blake;
-   keccak256_8way_context    keccak;
+   keccak256_4way_context    keccak;
    cubehashParam             cube;
-   skein256_8way_context     skein;
+   skein256_4way_context     skein;
    bmw256_8way_context       bmw;
 } lyra2v2_8way_ctx_holder __attribute__ ((aligned (64)));
 
@@ -116,9 +225,9 @@ static lyra2v2_8way_ctx_holder l2v2_8way_ctx;
 
 bool init_lyra2rev2_8way_ctx()
 {
-   keccak256_8way_init( &l2v2_8way_ctx.keccak );
+   keccak256_4way_init( &l2v2_8way_ctx.keccak );
    cubehashInit( &l2v2_8way_ctx.cube, 256, 16, 32 );
-   skein256_8way_init( &l2v2_8way_ctx.skein );
+   skein256_4way_init( &l2v2_8way_ctx.skein );
    bmw256_8way_init( &l2v2_8way_ctx.bmw );
    return true;
 }
@@ -126,7 +235,6 @@ bool init_lyra2rev2_8way_ctx()
 void lyra2rev2_8way_hash( void *state, const void *input )
 {
    uint32_t vhash[8*8] __attribute__ ((aligned (128)));
-   uint32_t vhashA[8*8] __attribute__ ((aligned (64)));
    uint32_t hash0[8] __attribute__ ((aligned (64)));
    uint32_t hash1[8] __attribute__ ((aligned (64)));
    uint32_t hash2[8] __attribute__ ((aligned (64)));
@@ -141,14 +249,19 @@ void lyra2rev2_8way_hash( void *state, const void *input )
    blake256_8way_update( &ctx.blake, input + (64<<3), 16 );
    blake256_8way_close( &ctx.blake, vhash );
 
-   rintrlv_8x32_8x64( vhashA, vhash, 256 );
-
-   keccak256_8way_update( &ctx.keccak, vhashA, 32 );
-   keccak256_8way_close( &ctx.keccak, vhash );
-
-   dintrlv_8x64( hash0, hash1, hash2, hash3,
+   dintrlv_8x32( hash0, hash1, hash2, hash3,
                  hash4, hash5, hash6, hash7, vhash, 256 );
 
+   intrlv_4x64( vhash, hash0, hash1, hash2, hash3, 256 );
+   keccak256_4way_update( &ctx.keccak, vhash, 32 );
+   keccak256_4way_close( &ctx.keccak, vhash );
+   dintrlv_4x64( hash0, hash1, hash2, hash3, vhash, 256 );
+   intrlv_4x64( vhash, hash4, hash5, hash6, hash7, 256 );
+   keccak256_4way_init( &ctx.keccak );
+   keccak256_4way_update( &ctx.keccak, vhash, 32 );
+   keccak256_4way_close( &ctx.keccak, vhash );
+   dintrlv_4x64( hash4, hash5, hash6, hash7, vhash, 256 );
+
    cubehash_full( &ctx.cube, (byte*) hash0, 256, (const byte*) hash0, 32 );
    cubehash_full( &ctx.cube, (byte*) hash1, 256, (const byte*) hash1, 32 );
    cubehash_full( &ctx.cube, (byte*) hash2, 256, (const byte*) hash2, 32 );
@@ -158,27 +271,25 @@ void lyra2rev2_8way_hash( void *state, const void *input )
    cubehash_full( &ctx.cube, (byte*) hash6, 256, (const byte*) hash6, 32 );
    cubehash_full( &ctx.cube, (byte*) hash7, 256, (const byte*) hash7, 32 );
 
-   intrlv_2x256( vhash, hash0, hash1, 256 );
-   LYRA2REV2_2WAY( l2v2_wholeMatrix, vhash, 32, vhash, 32, 1, 4, 4 );
-   dintrlv_2x256( hash0, hash1, vhash, 256 );
-   intrlv_2x256( vhash, hash2, hash3, 256 );
-   LYRA2REV2_2WAY( l2v2_wholeMatrix, vhash, 32, vhash, 32, 1, 4, 4 );
-   dintrlv_2x256( hash2, hash3, vhash, 256 );
-   intrlv_2x256( vhash, hash4, hash5, 256 );
-   LYRA2REV2_2WAY( l2v2_wholeMatrix, vhash, 32, vhash, 32, 1, 4, 4 );
-   dintrlv_2x256( hash4, hash5, vhash, 256 );
-   intrlv_2x256( vhash, hash6, hash7, 256 );
-   LYRA2REV2_2WAY( l2v2_wholeMatrix, vhash, 32, vhash, 32, 1, 4, 4 );
-   dintrlv_2x256( hash6, hash7, vhash, 256 );
 
-   intrlv_8x64( vhash, hash0, hash1, hash2, hash3, hash4, hash5, hash6,
-                hash7, 256 );
+   LYRA2REV2( l2v2_wholeMatrix, hash0, 32, hash0, 32, hash0, 32, 1, 4, 4 );
+   LYRA2REV2( l2v2_wholeMatrix, hash1, 32, hash1, 32, hash1, 32, 1, 4, 4 );
+   LYRA2REV2( l2v2_wholeMatrix, hash2, 32, hash2, 32, hash2, 32, 1, 4, 4 );
+   LYRA2REV2( l2v2_wholeMatrix, hash3, 32, hash3, 32, hash3, 32, 1, 4, 4 );
+   LYRA2REV2( l2v2_wholeMatrix, hash4, 32, hash4, 32, hash4, 32, 1, 4, 4 );
+   LYRA2REV2( l2v2_wholeMatrix, hash5, 32, hash5, 32, hash5, 32, 1, 4, 4 );
+   LYRA2REV2( l2v2_wholeMatrix, hash6, 32, hash6, 32, hash6, 32, 1, 4, 4 );
+   LYRA2REV2( l2v2_wholeMatrix, hash7, 32, hash7, 32, hash7, 32, 1, 4, 4 );
    
-   skein256_8way_update( &ctx.skein, vhash, 32 );
-   skein256_8way_close( &ctx.skein, vhash );
-
-   dintrlv_8x64( hash0, hash1, hash2, hash3,
-                 hash4, hash5, hash6, hash7, vhash, 256 );
+   intrlv_4x64( vhash, hash0, hash1, hash2, hash3, 256 );
+   skein256_4way_update( &ctx.skein, vhash, 32 );
+   skein256_4way_close( &ctx.skein, vhash );
+   dintrlv_4x64( hash0, hash1, hash2, hash3, vhash, 256 );
+   intrlv_4x64( vhash, hash4, hash5, hash6, hash7, 256 );
+   skein256_4way_init( &ctx.skein );
+   skein256_4way_update( &ctx.skein, vhash, 32 );
+   skein256_4way_close( &ctx.skein, vhash );
+   dintrlv_4x64( hash4, hash5, hash6, hash7, vhash, 256 );
 
    cubehash_full( &ctx.cube, (byte*) hash0, 256, (const byte*) hash0, 32 );
    cubehash_full( &ctx.cube, (byte*) hash1, 256, (const byte*) hash1, 32 );
@@ -189,8 +300,8 @@ void lyra2rev2_8way_hash( void *state, const void *input )
    cubehash_full( &ctx.cube, (byte*) hash6, 256, (const byte*) hash6, 32 );
    cubehash_full( &ctx.cube, (byte*) hash7, 256, (const byte*) hash7, 32 );
 
-   intrlv_8x32( vhash, hash0, hash1, hash2, hash3, hash4, hash5, hash6, 
-                hash7, 256 );
+   intrlv_8x32( vhash, hash0, hash1, hash2, hash3,
+                       hash4, hash5, hash6, hash7, 256 );
 
    bmw256_8way_update( &ctx.bmw, vhash, 32 );
    bmw256_8way_close( &ctx.bmw, state );
@@ -223,7 +334,6 @@ int scanhash_lyra2rev2_8way( struct work *work, const uint32_t max_nonce,
    do
    {
       lyra2rev2_8way_hash( hash, vdata );
-      pdata[19] = n;
 
       for ( int lane = 0; lane < 8; lane++ )
       if ( unlikely( hashd7[lane] <= targ32 ) )
@@ -243,6 +353,9 @@ int scanhash_lyra2rev2_8way( struct work *work, const uint32_t max_nonce,
    return 0;
 }
 
+#endif
+
+/*
 #elif defined (LYRA2REV2_4WAY)
 
 typedef struct {
@@ -367,3 +480,4 @@ int scanhash_lyra2rev2_4way( struct work *work, uint32_t max_nonce,
 }
 
 #endif
+*/

algo/scrypt/scrypt.c

@@ -380,7 +380,7 @@ static inline void PBKDF2_SHA256_128_32_8way(uint32_t *tstate,
 #endif /* HAVE_SHA256_8WAY */
 
 
-#if defined(USE_ASM) && defined(__x86_64__)
+//#if defined(USE_ASM) && defined(__x86_64__)
 
 #define SCRYPT_MAX_WAYS 12
 #define HAVE_SCRYPT_3WAY 1
@@ -394,113 +394,6 @@ void scrypt_core_3way(uint32_t *X, uint32_t *V, int N);
 void scrypt_core_6way(uint32_t *X, uint32_t *V, int N);
 #endif
 
-#elif defined(USE_ASM) && defined(__i386__)
-
-#define SCRYPT_MAX_WAYS 4
-#define scrypt_best_throughput() 1
-void scrypt_core(uint32_t *X, uint32_t *V, int N);
-
-#elif defined(USE_ASM) && defined(__arm__) && defined(__APCS_32__)
-
-void scrypt_core(uint32_t *X, uint32_t *V, int N);
-#if defined(__ARM_NEON__)
-#undef HAVE_SHA256_4WAY
-#define SCRYPT_MAX_WAYS 3
-#define HAVE_SCRYPT_3WAY 1
-#define scrypt_best_throughput() 3
-void scrypt_core_3way(uint32_t *X, uint32_t *V, int N);
-#endif
-
-#else
-
-static inline void xor_salsa8(uint32_t B[16], const uint32_t Bx[16])
-{
-	uint32_t x00,x01,x02,x03,x04,x05,x06,x07,x08,x09,x10,x11,x12,x13,x14,x15;
-	int i;
-
-	x00 = (B[ 0] ^= Bx[ 0]);
-	x01 = (B[ 1] ^= Bx[ 1]);
-	x02 = (B[ 2] ^= Bx[ 2]);
-	x03 = (B[ 3] ^= Bx[ 3]);
-	x04 = (B[ 4] ^= Bx[ 4]);
-	x05 = (B[ 5] ^= Bx[ 5]);
-	x06 = (B[ 6] ^= Bx[ 6]);
-	x07 = (B[ 7] ^= Bx[ 7]);
-	x08 = (B[ 8] ^= Bx[ 8]);
-	x09 = (B[ 9] ^= Bx[ 9]);
-	x10 = (B[10] ^= Bx[10]);
-	x11 = (B[11] ^= Bx[11]);
-	x12 = (B[12] ^= Bx[12]);
-	x13 = (B[13] ^= Bx[13]);
-	x14 = (B[14] ^= Bx[14]);
-	x15 = (B[15] ^= Bx[15]);
-	for (i = 0; i < 8; i += 2) {
-#define R(a, b) (((a) << (b)) | ((a) >> (32 - (b))))
-		/* Operate on columns. */
-		x04 ^= R(x00+x12, 7);	x09 ^= R(x05+x01, 7);
-		x14 ^= R(x10+x06, 7);	x03 ^= R(x15+x11, 7);
-		
-		x08 ^= R(x04+x00, 9);	x13 ^= R(x09+x05, 9);
-		x02 ^= R(x14+x10, 9);	x07 ^= R(x03+x15, 9);
-		
-		x12 ^= R(x08+x04,13);	x01 ^= R(x13+x09,13);
-		x06 ^= R(x02+x14,13);	x11 ^= R(x07+x03,13);
-		
-		x00 ^= R(x12+x08,18);	x05 ^= R(x01+x13,18);
-		x10 ^= R(x06+x02,18);	x15 ^= R(x11+x07,18);
-		
-		/* Operate on rows. */
-		x01 ^= R(x00+x03, 7);	x06 ^= R(x05+x04, 7);
-		x11 ^= R(x10+x09, 7);	x12 ^= R(x15+x14, 7);
-		
-		x02 ^= R(x01+x00, 9);	x07 ^= R(x06+x05, 9);
-		x08 ^= R(x11+x10, 9);	x13 ^= R(x12+x15, 9);
-		
-		x03 ^= R(x02+x01,13);	x04 ^= R(x07+x06,13);
-		x09 ^= R(x08+x11,13);	x14 ^= R(x13+x12,13);
-		
-		x00 ^= R(x03+x02,18);	x05 ^= R(x04+x07,18);
-		x10 ^= R(x09+x08,18);	x15 ^= R(x14+x13,18);
-#undef R
-	}
-	B[ 0] += x00;
-	B[ 1] += x01;
-	B[ 2] += x02;
-	B[ 3] += x03;
-	B[ 4] += x04;
-	B[ 5] += x05;
-	B[ 6] += x06;
-	B[ 7] += x07;
-	B[ 8] += x08;
-	B[ 9] += x09;
-	B[10] += x10;
-	B[11] += x11;
-	B[12] += x12;
-	B[13] += x13;
-	B[14] += x14;
-	B[15] += x15;
-}
-
-static inline void scrypt_core(uint32_t *X, uint32_t *V, int N)
-{
-	int i;
-
-	for (i = 0; i < N; i++) {
-		memcpy(&V[i * 32], X, 128);
-		xor_salsa8(&X[0], &X[16]);
-		xor_salsa8(&X[16], &X[0]);
-	}
-	for (i = 0; i < N; i++) {
-		uint32_t j = 32 * (X[16] & (N - 1));
-		for (uint8_t k = 0; k < 32; k++)
-			X[k] ^= V[j + k];
-		xor_salsa8(&X[0], &X[16]);
-		xor_salsa8(&X[16], &X[0]);
-	}
-}
-
-#endif
-
 #ifndef SCRYPT_MAX_WAYS
 #define SCRYPT_MAX_WAYS 1
 #define scrypt_best_throughput() 1
@@ -511,8 +404,8 @@ unsigned char *scrypt_buffer_alloc(int N)
 	return (uchar*) malloc((size_t)N * SCRYPT_MAX_WAYS * 128 + 63);
 }
 
-static void scrypt_1024_1_1_256(const uint32_t *input, uint32_t *output,
-	uint32_t *midstate, unsigned char *scratchpad, int N)
+static bool scrypt_1024_1_1_256(const uint32_t *input, uint32_t *output,
+	uint32_t *midstate, unsigned char *scratchpad, int N, int thr_id )
 {
 	uint32_t tstate[8], ostate[8];
 	uint32_t X[32];
@@ -527,11 +420,13 @@ static void scrypt_1024_1_1_256(const uint32_t *input, uint32_t *output,
 	scrypt_core(X, V, N);
 
 	PBKDF2_SHA256_128_32(tstate, ostate, X, output);
+   return true;
 }
 
 #ifdef HAVE_SHA256_4WAY
-static void scrypt_1024_1_1_256_4way(const uint32_t *input,
-	uint32_t *output, uint32_t *midstate, unsigned char *scratchpad, int N)
+static bool scrypt_1024_1_1_256_4way(const uint32_t *input,
+	uint32_t *output, uint32_t *midstate, unsigned char *scratchpad, int N, 
+   int thrid )
 {
 	uint32_t _ALIGN(128) tstate[4 * 8];
 	uint32_t _ALIGN(128) ostate[4 * 8];
@@ -545,32 +440,43 @@ static void scrypt_1024_1_1_256_4way(const uint32_t *input,
 	for (i = 0; i < 20; i++)
 		for (k = 0; k < 4; k++)
 			W[4 * i + k] = input[k * 20 + i];
+
    for (i = 0; i < 8; i++)
 		for (k = 0; k < 4; k++)
 			tstate[4 * i + k] = midstate[i];
+
    HMAC_SHA256_80_init_4way(W, tstate, ostate);
+
    PBKDF2_SHA256_80_128_4way(tstate, ostate, W, W);
+
    for (i = 0; i < 32; i++)
 		for (k = 0; k < 4; k++)
 			X[k * 32 + i] = W[4 * i + k];
+
    scrypt_core(X + 0 * 32, V, N);
 	scrypt_core(X + 1 * 32, V, N);
 	scrypt_core(X + 2 * 32, V, N);
 	scrypt_core(X + 3 * 32, V, N);
+
    for (i = 0; i < 32; i++)
 		for (k = 0; k < 4; k++)
 			W[4 * i + k] = X[k * 32 + i];
+
    PBKDF2_SHA256_128_32_4way(tstate, ostate, W, W);
+
    for (i = 0; i < 8; i++)
 		for (k = 0; k < 4; k++)
 			output[k * 8 + i] = W[4 * i + k];
+
+   return true;
 }
 #endif /* HAVE_SHA256_4WAY */
 
 #ifdef HAVE_SCRYPT_3WAY
 
-static void scrypt_1024_1_1_256_3way(const uint32_t *input,
-	uint32_t *output, uint32_t *midstate, unsigned char *scratchpad, int N)
+static bool scrypt_1024_1_1_256_3way(const uint32_t *input,
+	uint32_t *output, uint32_t *midstate, unsigned char *scratchpad, int N, 
+   int thrid )
 {
 	uint32_t _ALIGN(64) tstate[3 * 8], ostate[3 * 8];
 	uint32_t _ALIGN(64) X[3 * 32];
@@ -581,23 +487,34 @@ static void scrypt_1024_1_1_256_3way(const uint32_t *input,
 	memcpy(tstate +  0, midstate, 32);
 	memcpy(tstate +  8, midstate, 32);
 	memcpy(tstate + 16, midstate, 32);
+
    HMAC_SHA256_80_init(input +  0, tstate +  0, ostate +  0);
 	HMAC_SHA256_80_init(input + 20, tstate +  8, ostate +  8);
 	HMAC_SHA256_80_init(input + 40, tstate + 16, ostate + 16);
+
+   if ( work_restart[thrid].restart ) return false;
+
    PBKDF2_SHA256_80_128(tstate +  0, ostate +  0, input +  0, X +  0);
 	PBKDF2_SHA256_80_128(tstate +  8, ostate +  8, input + 20, X + 32);
 	PBKDF2_SHA256_80_128(tstate + 16, ostate + 16, input + 40, X + 64);
 
+   if ( work_restart[thrid].restart ) return false;
+
    scrypt_core_3way(X, V, N);
 
+   if ( work_restart[thrid].restart ) return false;
+
    PBKDF2_SHA256_128_32(tstate +  0, ostate +  0, X +  0, output +  0);
 	PBKDF2_SHA256_128_32(tstate +  8, ostate +  8, X + 32, output +  8);
 	PBKDF2_SHA256_128_32(tstate + 16, ostate + 16, X + 64, output + 16);
+
+   return true;
 }
 
 #ifdef HAVE_SHA256_4WAY
-static void scrypt_1024_1_1_256_12way(const uint32_t *input,
-	uint32_t *output, uint32_t *midstate, unsigned char *scratchpad, int N)
+static bool scrypt_1024_1_1_256_12way(const uint32_t *input,
+	uint32_t *output, uint32_t *midstate, unsigned char *scratchpad, int N,
+   int thrid )
 {
 	uint32_t _ALIGN(128) tstate[12 * 8];
 	uint32_t _ALIGN(128) ostate[12 * 8];
@@ -612,43 +529,60 @@ static void scrypt_1024_1_1_256_12way(const uint32_t *input,
 		for (i = 0; i < 20; i++)
 			for (k = 0; k < 4; k++)
 				W[128 * j + 4 * i + k] = input[80 * j + k * 20 + i];
+
    for (j = 0; j < 3; j++)
 		for (i = 0; i < 8; i++)
 			for (k = 0; k < 4; k++)
 				tstate[32 * j + 4 * i + k] = midstate[i];
+
    HMAC_SHA256_80_init_4way(W +   0, tstate +  0, ostate +  0);
 	HMAC_SHA256_80_init_4way(W + 128, tstate + 32, ostate + 32);
 	HMAC_SHA256_80_init_4way(W + 256, tstate + 64, ostate + 64);
+
+   if ( work_restart[thrid].restart ) return false;
+
    PBKDF2_SHA256_80_128_4way(tstate +  0, ostate +  0, W +   0, W +   0);
 	PBKDF2_SHA256_80_128_4way(tstate + 32, ostate + 32, W + 128, W + 128);
 	PBKDF2_SHA256_80_128_4way(tstate + 64, ostate + 64, W + 256, W + 256);
+
+   if ( work_restart[thrid].restart ) return false;
+
    for (j = 0; j < 3; j++)
 		for (i = 0; i < 32; i++)
 			for (k = 0; k < 4; k++)
 				X[128 * j + k * 32 + i] = W[128 * j + 4 * i + k];
+
    scrypt_core_3way(X + 0 * 96, V, N);
 	scrypt_core_3way(X + 1 * 96, V, N);
 	scrypt_core_3way(X + 2 * 96, V, N);
 	scrypt_core_3way(X + 3 * 96, V, N);
+
+   if ( work_restart[thrid].restart ) return false;
+
    for (j = 0; j < 3; j++)
 		for (i = 0; i < 32; i++)
 			for (k = 0; k < 4; k++)
 				W[128 * j + 4 * i + k] = X[128 * j + k * 32 + i];
+
    PBKDF2_SHA256_128_32_4way(tstate +  0, ostate +  0, W +   0, W +   0);
 	PBKDF2_SHA256_128_32_4way(tstate + 32, ostate + 32, W + 128, W + 128);
 	PBKDF2_SHA256_128_32_4way(tstate + 64, ostate + 64, W + 256, W + 256);
+
    for (j = 0; j < 3; j++)
 		for (i = 0; i < 8; i++)
 			for (k = 0; k < 4; k++)
 				output[32 * j + k * 8 + i] = W[128 * j + 4 * i + k];
+
+   return true;
 }
 #endif /* HAVE_SHA256_4WAY */
 
 #endif /* HAVE_SCRYPT_3WAY */
 
 #ifdef HAVE_SCRYPT_6WAY
-static void scrypt_1024_1_1_256_24way(const uint32_t *input,
-	uint32_t *output, uint32_t *midstate, unsigned char *scratchpad, int N)
+static bool scrypt_1024_1_1_256_24way( const uint32_t *input,
+                               uint32_t *output, uint32_t *midstate,
+                               unsigned char *scratchpad, int N, int thrid )
 {
 	uint32_t _ALIGN(128) tstate[24 * 8];
 	uint32_t _ALIGN(128) ostate[24 * 8];
@@ -663,35 +597,51 @@ static void scrypt_1024_1_1_256_24way(const uint32_t *input,
 		for ( i = 0; i < 20; i++ )
 			for ( k = 0; k < 8; k++ )
 				W[8 * 32 * j + 8 * i + k] = input[8 * 20 * j + k * 20 + i];
+
    for ( j = 0; j < 3; j++ )
 		for ( i = 0; i < 8; i++ )
 			for ( k = 0; k < 8; k++ )
 				tstate[8 * 8 * j + 8 * i + k] = midstate[i];
+
    HMAC_SHA256_80_init_8way( W +   0, tstate +   0, ostate +   0 );
 	HMAC_SHA256_80_init_8way( W + 256, tstate +  64, ostate +  64 );
 	HMAC_SHA256_80_init_8way( W + 512, tstate + 128, ostate + 128 );
+
+   if ( work_restart[thrid].restart ) return false;
+   
    PBKDF2_SHA256_80_128_8way( tstate +   0, ostate +   0, W +   0, W +   0 );
 	PBKDF2_SHA256_80_128_8way( tstate +  64, ostate +  64, W + 256, W + 256 );
 	PBKDF2_SHA256_80_128_8way( tstate + 128, ostate + 128, W + 512, W + 512 );
+
+   if ( work_restart[thrid].restart ) return false;
+
    for ( j = 0; j < 3; j++ )
 		for ( i = 0; i < 32; i++ )
 			for ( k = 0; k < 8; k++ )
 				X[8 * 32 * j + k * 32 + i] = W[8 * 32 * j + 8 * i + k];
+
    scrypt_core_6way( X +  0 * 32, V, N );
 	scrypt_core_6way( X +  6 * 32, V, N );
 	scrypt_core_6way( X + 12 * 32, V, N );
 	scrypt_core_6way( X + 18 * 32, V, N );
+
+   if ( work_restart[thrid].restart ) return false;
+   
    for ( j = 0; j < 3; j++ )
 		for ( i = 0; i < 32; i++ )
 			for ( k = 0; k < 8; k++ )
 				W[8 * 32 * j + 8 * i + k] = X[8 * 32 * j + k * 32 + i];
+
    PBKDF2_SHA256_128_32_8way( tstate +   0, ostate +   0, W +   0, W +   0 );
 	PBKDF2_SHA256_128_32_8way( tstate +  64, ostate +  64, W + 256, W + 256 );
 	PBKDF2_SHA256_128_32_8way( tstate + 128, ostate + 128, W + 512, W + 512 );
+
    for ( j = 0; j < 3; j++ )
 		for ( i = 0; i < 8; i++ )
 			for ( k = 0; k < 8; k++ )
 				output[8 * 8 * j + k * 8 + i] = W[8 * 32 * j + 8 * i + k];
+
+   return true;
 }
 #endif /* HAVE_SCRYPT_6WAY */
 
@@ -703,7 +653,6 @@ extern int scanhash_scrypt( struct work *work, uint32_t max_nonce,
 	uint32_t data[SCRYPT_MAX_WAYS * 20], hash[SCRYPT_MAX_WAYS * 8];
 	uint32_t midstate[8];
 	uint32_t n = pdata[19] - 1;
-	const uint32_t Htarg = ptarget[7];
    int thr_id = mythr->id;  // thr_id arg is deprecated
 	int throughput = scrypt_best_throughput();
 	int i;
@@ -714,6 +663,8 @@ extern int scanhash_scrypt( struct work *work, uint32_t max_nonce,
 		throughput *= 4;
 #endif
 	
+//   applog(LOG_INFO,"Scrypt thoughput %d",throughput);
+
 	for (i = 0; i < throughput; i++)
 		memcpy(data + i * 20, pdata, 80);
 	
@@ -721,46 +672,49 @@ extern int scanhash_scrypt( struct work *work, uint32_t max_nonce,
 	sha256_transform(midstate, data, 0);
 	
 	do {
-
+      bool rc = true;
 		for (i = 0; i < throughput; i++)
 			data[i * 20 + 19] = ++n;
 		
 #if defined(HAVE_SHA256_4WAY)
 		if (throughput == 4)
-			scrypt_1024_1_1_256_4way(data, hash, midstate,
-                             scratchbuf, scratchbuf_size );
+			rc = scrypt_1024_1_1_256_4way(data, hash, midstate,
+                             scratchbuf, scratchbuf_size, thr_id );
 		else
 #endif
 #if defined(HAVE_SCRYPT_3WAY) && defined(HAVE_SHA256_4WAY)
 		if (throughput == 12)
-			scrypt_1024_1_1_256_12way(data, hash, midstate,
-                              scratchbuf, scratchbuf_size );
+			rc = scrypt_1024_1_1_256_12way(data, hash, midstate,
+                              scratchbuf, scratchbuf_size, thr_id );
 		else
 #endif
 #if defined(HAVE_SCRYPT_6WAY)
 		if (throughput == 24)
-			scrypt_1024_1_1_256_24way(data, hash, midstate,
-                               scratchbuf, scratchbuf_size );
+			rc = scrypt_1024_1_1_256_24way(data, hash, midstate,
+                               scratchbuf, scratchbuf_size, thr_id );
 		else
 #endif
 #if defined(HAVE_SCRYPT_3WAY)
 		if (throughput == 3)
-			scrypt_1024_1_1_256_3way(data, hash, midstate,
-                                scratchbuf, scratchbuf_size );
+			rc = scrypt_1024_1_1_256_3way(data, hash, midstate,
+                                scratchbuf, scratchbuf_size, thr_id );
 		else
 #endif
-		scrypt_1024_1_1_256(data, hash, midstate, scratchbuf,
-                                scratchbuf_size );
+		rc = scrypt_1024_1_1_256(data, hash, midstate, scratchbuf,
+                                scratchbuf_size, thr_id );
 		
-		for (i = 0; i < throughput; i++) {
-			if (unlikely(hash[i * 8 + 7] <= Htarg && fulltest(hash + i * 8, ptarget))) {
+      if ( rc )
+      for ( i = 0; i < throughput; i++ )
+      {
+         if ( unlikely( valid_hash( hash + i * 8, ptarget ) ) )
+         {
 				pdata[19] = data[i * 20 + 19];
-            submit_solution( work, hash, mythr );
+            submit_solution( work, hash + i * 8, mythr );
 			}
       }
-	} while ( likely( n < max_nonce && !(*restart) ) );
+	} while ( likely( ( n < ( max_nonce - throughput ) ) && !(*restart) ) );
 	
-	*hashes_done = n - pdata[19] + 1;
+	*hashes_done = n - pdata[19];
 	pdata[19] = n;
 	return 0;
 }
@@ -779,7 +733,6 @@ bool register_scrypt_algo( algo_gate_t* gate )
   gate->optimizations = SSE2_OPT | AVX2_OPT;
   gate->miner_thread_init =(void*)&scrypt_miner_thread_init;
   gate->scanhash         = (void*)&scanhash_scrypt;
-//  gate->hash             = (void*)&scrypt_1024_1_1_256_24way;
   opt_target_factor = 65536.0;
 
   if ( !opt_param_n )

configure

@@ -1,6 +1,6 @@
 #! /bin/sh
 # Guess values for system-dependent variables and create Makefiles.
-# Generated by GNU Autoconf 2.69 for cpuminer-opt 3.12.3.1.
+# Generated by GNU Autoconf 2.69 for cpuminer-opt 3.12.5.
 #
 #
 # Copyright (C) 1992-1996, 1998-2012 Free Software Foundation, Inc.
@@ -577,8 +577,8 @@ MAKEFLAGS=
 # Identity of this package.
 PACKAGE_NAME='cpuminer-opt'
 PACKAGE_TARNAME='cpuminer-opt'
-PACKAGE_VERSION='3.12.3.1'
-PACKAGE_STRING='cpuminer-opt 3.12.3.1'
+PACKAGE_VERSION='3.12.5'
+PACKAGE_STRING='cpuminer-opt 3.12.5'
 PACKAGE_BUGREPORT=''
 PACKAGE_URL=''
 
@@ -1332,7 +1332,7 @@ if test "$ac_init_help" = "long"; then
   # Omit some internal or obsolete options to make the list less imposing.
   # This message is too long to be a string in the A/UX 3.1 sh.
   cat <<_ACEOF
-\`configure' configures cpuminer-opt 3.12.3.1 to adapt to many kinds of systems.
+\`configure' configures cpuminer-opt 3.12.5 to adapt to many kinds of systems.
 
 Usage: $0 [OPTION]... [VAR=VALUE]...
 
@@ -1404,7 +1404,7 @@ fi
 
 if test -n "$ac_init_help"; then
   case $ac_init_help in
-     short | recursive ) echo "Configuration of cpuminer-opt 3.12.3.1:";;
+     short | recursive ) echo "Configuration of cpuminer-opt 3.12.5:";;
    esac
   cat <<\_ACEOF
 
@@ -1509,7 +1509,7 @@ fi
 test -n "$ac_init_help" && exit $ac_status
 if $ac_init_version; then
   cat <<\_ACEOF
-cpuminer-opt configure 3.12.3.1
+cpuminer-opt configure 3.12.5
 generated by GNU Autoconf 2.69
 
 Copyright (C) 2012 Free Software Foundation, Inc.
@@ -2012,7 +2012,7 @@ cat >config.log <<_ACEOF
 This file contains any messages produced by compilers while
 running configure, to aid debugging if configure makes a mistake.
 
-It was created by cpuminer-opt $as_me 3.12.3.1, which was
+It was created by cpuminer-opt $as_me 3.12.5, which was
 generated by GNU Autoconf 2.69.  Invocation command line was
 
   $ $0 $@
@@ -2993,7 +2993,7 @@ fi
 
 # Define the identity of the package.
  PACKAGE='cpuminer-opt'
- VERSION='3.12.3.1'
+ VERSION='3.12.5'
 
 
 cat >>confdefs.h <<_ACEOF
@@ -6690,7 +6690,7 @@ cat >>$CONFIG_STATUS <<\_ACEOF || ac_write_fail=1
 # report actual input values of CONFIG_FILES etc. instead of their
 # values after options handling.
 ac_log="
-This file was extended by cpuminer-opt $as_me 3.12.3.1, which was
+This file was extended by cpuminer-opt $as_me 3.12.5, which was
 generated by GNU Autoconf 2.69.  Invocation command line was
 
   CONFIG_FILES    = $CONFIG_FILES
@@ -6756,7 +6756,7 @@ _ACEOF
 cat >>$CONFIG_STATUS <<_ACEOF || ac_write_fail=1
 ac_cs_config="`$as_echo "$ac_configure_args" | sed 's/^ //; s/[\\""\`\$]/\\\\&/g'`"
 ac_cs_version="\\
-cpuminer-opt config.status 3.12.3.1
+cpuminer-opt config.status 3.12.5
 configured by $0, generated by GNU Autoconf 2.69,
   with options \\"\$ac_cs_config\\"
 

configure.ac

@@ -1,4 +1,4 @@
-AC_INIT([cpuminer-opt], [3.12.3.1])
+AC_INIT([cpuminer-opt], [3.12.5])
 
 AC_PREREQ([2.59c])
 AC_CANONICAL_SYSTEM

cpu-miner.c

@@ -102,6 +102,7 @@ static int opt_fail_pause = 10;
 static int opt_time_limit = 0;
 int opt_timeout = 300;
 static int opt_scantime = 5;
+const int min_scantime = 1;
 //static const bool opt_time = true;
 enum algos opt_algo = ALGO_NULL;
 char* opt_param_key = NULL;
@@ -159,6 +160,7 @@ uint32_t accepted_share_count = 0;
 uint32_t rejected_share_count = 0;
 uint32_t stale_share_count = 0;
 uint32_t solved_block_count = 0;
+uint64_t new_work_count = 0;
 double *thr_hashrates;
 double global_hashrate = 0;
 double stratum_diff = 0.;
@@ -184,6 +186,22 @@ int default_api_listen = 4048;
   pthread_mutex_t applog_lock;
   pthread_mutex_t stats_lock;
 
+static struct   timeval session_start;
+static struct   timeval five_min_start;
+static uint64_t session_first_block = 0;
+static double   latency_sum = 0.;
+static uint64_t submit_sum  = 0;
+static uint64_t accept_sum  = 0;
+static uint64_t stale_sum  = 0;
+static uint64_t reject_sum  = 0;
+static double   norm_diff_sum = 0.;
+static uint32_t last_block_height = 0;
+//static bool     new_job = false;
+static double   last_targetdiff = 0.;
+#if !(defined(__WINDOWS__) || defined(_WIN64) || defined(_WIN32))
+static uint32_t hi_temp = 0;
+#endif
+
   
 static char const short_options[] =
 #ifdef HAVE_SYSLOG_H
@@ -420,8 +438,7 @@ static bool work_decode( const json_t *val, struct work *work )
     if ( !allow_mininginfo )
         net_diff = algo_gate.calc_network_diff( work );
     work->targetdiff = target_to_diff( work->target );
-    // for api stats, on longpoll pools
-    stratum_diff = work->targetdiff;
+    stratum_diff = last_targetdiff = work->targetdiff;
     work->sharediff = 0;
     algo_gate.decode_extra_data( work, &net_blocks );
     return true;
@@ -437,17 +454,17 @@ static bool get_mininginfo(CURL *curl, struct work *work)
 		return false;
 
 	int curl_err = 0;
-	json_t *val = json_rpc_call(curl, rpc_url, rpc_userpass, info_req, &curl_err, 0);
+	json_t *val = json_rpc_call( curl, rpc_url, rpc_userpass, info_req,
+                                &curl_err, 0 );
 
-	if (!val && curl_err == -1) {
+	if ( !val && curl_err == -1 )
+   {
 		allow_mininginfo = false;
-		if (opt_debug) {
+		if ( opt_debug )
 			applog( LOG_DEBUG, "getmininginfo not supported" );
-		}
 		return false;
 	}
-	else
-        {
+
    json_t *res = json_object_get( val, "result" );
    // "blocks": 491493 (= current work height - 1)
    // "difficulty": 0.99607860999999998
@@ -455,20 +472,26 @@ static bool get_mininginfo(CURL *curl, struct work *work)
    if ( res )
    {
   		json_t *key = json_object_get( res, "difficulty" );
-		if (key) {
+   	if ( key )
+      {
 	   	if ( json_is_object( key ) )
 		   	key = json_object_get( key, "proof-of-work" );
 		   if ( json_is_real( key ) )
 			   net_diff = json_real_value( key );
 	   }
+
       key = json_object_get( res, "networkhashps" );
-		if (key && json_is_integer(key)) {
+	   if ( key && json_is_integer( key ) )
 		   net_hashrate = (double) json_integer_value( key );
-		}
+
       key = json_object_get( res, "blocks" );
-		if (key && json_is_integer(key)) {
+	   if ( key && json_is_integer( key ) )
 		  	net_blocks = json_integer_value( key );
-		}
+
+      if ( opt_debug )
+         applog(LOG_INFO," Mining info: diff %.5g, net_hashrate %f, height %d",
+                              net_diff, net_hashrate, net_blocks );
+      
       if ( !work->height )
       {
 	      // complete missing data from getwork
@@ -476,11 +499,15 @@ static bool get_mininginfo(CURL *curl, struct work *work)
 	      if ( work->height > g_work.height )
          {
             restart_threads();
-			if (!opt_quiet) {
+
+/* redundant with new block log
+            if ( !opt_quiet )
+            {
 		         char netinfo[64] = { 0 };
 		         char srate[32] = { 0 };
 		         sprintf( netinfo, "diff %.2f", net_diff );
-			   if (net_hashrate) {
+		         if ( net_hashrate )
+               {
 	               format_hashrate( net_hashrate, srate );
                   strcat( netinfo, ", net " );
 			         strcat( netinfo, srate );
@@ -488,9 +515,9 @@ static bool get_mininginfo(CURL *curl, struct work *work)
 		         applog( LOG_BLUE, "%s block %d, %s",
 			                algo_names[opt_algo], work->height, netinfo );
 		      }
+*/
          } 
-		}
-	   }
+	   }  // res
 	}
 	json_decref( val );
 	return true;
@@ -866,33 +893,8 @@ static inline void sprintf_et( char *str, int seconds )
       sprintf( str, "%um%02us", min, sec );
 }
    
-// Bitcoin formula for converting difficulty to an equivalent
-// number of hashes.
-//
-//     https://en.bitcoin.it/wiki/Difficulty
-//
-//     hash = diff * 2**32
-//
-// diff_to_hash = 2**32 = 0x100000000 = 4294967296;
-
-const double diff_to_hash = 4294967296.;
-
-static struct   timeval session_start;
-static struct   timeval five_min_start;
-static uint64_t session_first_block = 0;
-static double   latency_sum = 0.;
-static uint64_t submit_sum  = 0;
-static uint64_t accept_sum  = 0;
-static uint64_t stale_sum  = 0;
-static uint64_t reject_sum  = 0;
-static double   norm_diff_sum = 0.;
-static uint32_t last_block_height = 0;
-//static bool     new_job = false;
-static double   last_targetdiff = 0.;
-#if !(defined(__WINDOWS__) || defined(_WIN64) || defined(_WIN32))
-static uint32_t hi_temp = 0;
-#endif
-//static uint32_t stratum_errors = 0;
+const double exp32 = 4294967296.;                // 2**32
+const double exp64 = 4294967296. * 4294967296.;  // 2**64
 
 struct share_stats_t
 {
@@ -936,6 +938,8 @@ void report_summary_log( bool force )
    uint64_t accepts = accept_sum;  accept_sum = 0;
    uint64_t rejects = reject_sum;  reject_sum = 0;
    uint64_t stales  = stale_sum;   stale_sum  = 0;
+//   uint64_t new_work = new_work_count; new_work_count = 0;
+
    memcpy( &start_time, &five_min_start, sizeof start_time );
    memcpy( &five_min_start, &now, sizeof now );
 
@@ -946,12 +950,10 @@ void report_summary_log( bool force )
    
    double share_time = (double)et.tv_sec + (double)et.tv_usec / 1e6;
    double ghrate = global_hashrate;
-
-   double shrate = share_time == 0. ? 0. : diff_to_hash * last_targetdiff
+   double shrate = share_time == 0. ? 0. : exp32 * last_targetdiff
                                            * (double)(accepts) / share_time;
-   double sess_hrate = uptime.tv_sec == 0. ? 0. : diff_to_hash * norm_diff_sum
+   double sess_hrate = uptime.tv_sec == 0. ? 0. : exp32 * norm_diff_sum
                                                    / (double)uptime.tv_sec;
-
    double submit_rate = share_time == 0. ? 0. : (double)submits*60. / share_time;
    char shr_units[4] = {0};
    char ghr_units[4] = {0};
@@ -959,6 +961,21 @@ void report_summary_log( bool force )
    char et_str[24];
    char upt_str[24];
 
+/*
+   uint64_t work_time = new_work ? et.tv_sec / new_work : 15;
+   // Adjust scantime for frequent new work to prevent stales.
+   if ( work_time < 2 * opt_scantime )
+   {
+      if ( opt_scantime > min_scantime )
+      {
+         opt_scantime--;
+         if ( !opt_quiet )
+            applog( LOG_BLUE, "Getwork scan time reduced to %d seconds",
+                               opt_scantime );
+      }
+   }   
+*/
+
    scale_hash_for_display( &shrate, shr_units );
    scale_hash_for_display( &ghrate, ghr_units );
    scale_hash_for_display( &sess_hrate, sess_hr_units );
@@ -966,7 +983,8 @@ void report_summary_log( bool force )
    sprintf_et( et_str, et.tv_sec );
    sprintf_et( upt_str, uptime.tv_sec );
 
-   applog( LOG_NOTICE, "Periodic Report     %s        %s", et_str, upt_str );
+   applog( LOG_BLUE, "%s: %s", algo_names[ opt_algo ], short_url );
+   applog2( LOG_NOTICE, "Periodic Report     %s        %s", et_str, upt_str );
    applog2( LOG_INFO, "Share rate        %.2f/min     %.2f/min",
                       submit_rate, (double)submitted_share_count*60. /
                     ( (double)uptime.tv_sec + (double)uptime.tv_usec / 1e6 ) );
@@ -977,11 +995,11 @@ void report_summary_log( bool force )
    if ( accepted_share_count < submitted_share_count )
    {
       double lost_ghrate = uptime.tv_sec == 0. ? 0.
-                  : diff_to_hash * last_targetdiff
+                  : exp32 * last_targetdiff
                       * (double)(submitted_share_count - accepted_share_count )
                     / (double)uptime.tv_sec;
       double lost_shrate = share_time == 0. ? 0.
-               : diff_to_hash * last_targetdiff  * (double)(submits - accepts )
+               : exp32 * last_targetdiff  * (double)(submits - accepts )
                 / share_time;
       char lshr_units[4] = {0};
       char lghr_units[4] = {0};
@@ -1008,7 +1026,7 @@ void report_summary_log( bool force )
 
 bool lowdiff_debug = false;
 
-static int share_result( int result, struct work *null_work,
+static int share_result( int result, struct work *work,
                          const char *reason )
 {
    double share_time = 0., share_ratio = 0.;
@@ -1053,31 +1071,45 @@ static int share_result( int result, struct work *null_work,
    }
 
    share_ratio = my_stats.net_diff == 0. ? 0. : my_stats.share_diff /
-                                                my_stats.net_diff * 100.;
-
+                                                my_stats.net_diff;
    // check result
    if ( likely( result ) )
    {
       accepted_share_count++;
+      sprintf( sres, "S%d", stale_share_count );
+      sprintf( rres, "R%d", rejected_share_count );
       if unlikely( ( my_stats.net_diff > 0. )
                 && ( my_stats.share_diff >= net_diff ) )
       {
          solved = true;
          solved_block_count++;
+         sprintf( bres, "BLOCK SOLVED %d", solved_block_count );
+         sprintf( ares, "A%d", accepted_share_count );
+      }
+      else
+      {
+         sprintf( bres, "B%d", solved_block_count );
+         sprintf( ares, "Accepted %d", accepted_share_count );
       }
    }
    else
    {
-     if ( reason && strstr( reason, "Invalid job id" ) )
+     sprintf( ares, "A%d", accepted_share_count );
+     sprintf( bres, "B%d", solved_block_count );
+     if ( work ) stale = work->stale;
+     if ( reason ) stale = stale || strstr( reason, "Invalid job id" );
+     if ( stale )
      {
-        stale = true;
         stale_share_count++;
+        sprintf( sres, "Stale %d", stale_share_count );
+        sprintf( rres, "R%d", rejected_share_count );
      }
      else
      {
         rejected_share_count++;
-        lowdiff_debug = true;
-  
+        sprintf( sres, "S%d", stale_share_count );
+        sprintf( rres, "Rejected %d" , rejected_share_count );
+//        lowdiff_debug = true;
      }
    }
 
@@ -1103,59 +1135,34 @@ static int share_result( int result, struct work *null_work,
 
    pthread_mutex_unlock( &stats_lock );
 
-   if ( likely( result ) )
-   {
-     if ( unlikely( solved ) )
-     {
-       sprintf( bres, "BLOCK SOLVED %d", solved_block_count );
-       sprintf( ares, "A%d", accepted_share_count );
-     }
-     else
-     {
-       sprintf( bres, "B%d", solved_block_count );
-       sprintf( ares, "Accepted %d", accepted_share_count );
-     }
-     sprintf( sres, "S%d", stale_share_count );
-     sprintf( rres, "R%d", rejected_share_count );
-   }
-   else  
-   {
-     sprintf( ares, "A%d", accepted_share_count );
-     sprintf( bres, "B%d", solved_block_count );
-     if ( stale )
-     {
-        sprintf( sres, "Stale %d", stale_share_count );
-        sprintf( rres, "R%d", rejected_share_count );
-     }
-     else
-     {
-        sprintf( sres, "S%d", stale_share_count ); 
-        sprintf( rres, "Rejected %d" , rejected_share_count );
-     }
-   } 
-
    if ( use_colors )
    {
-     bcol = acol = scol = rcol = CL_WHT;
+     bcol = acol = scol = rcol = CL_N;
      if ( likely( result ) )
      {
-       acol = CL_GRN;  
-       if ( unlikely( solved ) ) bcol = CL_MAG;
+       acol = CL_WHT CL_GRN;  
+       if ( unlikely( solved ) ) bcol = CL_WHT CL_MAG;
      }        
-     else if ( stale ) scol = CL_YL2;
-     else              rcol = CL_RED;
+     else if ( stale ) scol = CL_WHT CL_YL2;
+     else              rcol = CL_WHT CL_RED;
    }
    else
       bcol = acol = scol = rcol = "\0";
 
-   applog( LOG_NOTICE, "%d %s%s %s%s %s%s %s%s" CL_WHT ", %.3f sec (%dms)",
+   applog( LOG_NOTICE, "%d %s%s %s%s %s%s %s%s" CL_N ", %.3f sec (%dms)",
            my_stats.share_count, acol, ares, scol, sres, rcol, rres, bcol,
            bres, share_time, latency );
 
-   if ( have_stratum && !opt_quiet )
-      applog2( LOG_NOTICE, "Diff %.5g (%.3g%), %sBlock %d, %sJob %s" CL_WHT,
+   if ( !opt_quiet )
+   {
+      if ( have_stratum )
+         applog2( LOG_INFO, "Diff %.5g (%.3g), %sBlock %d" CL_N ", %sJob %s",
                my_stats.share_diff, share_ratio, bcol, stratum.block_height,
                scol, my_stats.job_id );
+      else
+         applog2( LOG_INFO, "Diff %.5g (%.3g), %sBlock %d",
+               my_stats.share_diff, share_ratio, bcol, stratum.block_height );
+   }
 
    if ( unlikely( reason && !result ) )
    {
@@ -1188,6 +1195,9 @@ static int share_result( int result, struct work *null_work,
    return 1;
 }
 
+static const char *json_submit_req =
+   "{\"method\": \"mining.submit\", \"params\": [\"%s\", \"%s\", \"%s\", \"%s\", \"%s\"], \"id\":4}";
+
 void std_le_build_stratum_request( char *req, struct work *work )
 {
    unsigned char *xnonce2str;
@@ -1198,9 +1208,8 @@ void std_le_build_stratum_request( char *req, struct work *work )
    bin2hex( ntimestr, (char*)(&ntime), sizeof(uint32_t) );
    bin2hex( noncestr, (char*)(&nonce), sizeof(uint32_t) );
    xnonce2str = abin2hex( work->xnonce2, work->xnonce2_len );
-   snprintf( req, JSON_BUF_LEN,
-        "{\"method\": \"mining.submit\", \"params\": [\"%s\", \"%s\", \"%s\", \"%s\", \"%s\"], \"id\":4}",
-         rpc_user, work->job_id, xnonce2str, ntimestr, noncestr );
+   snprintf( req, JSON_BUF_LEN, json_submit_req, rpc_user, work->job_id,
+             xnonce2str, ntimestr, noncestr );
    free( xnonce2str );
 }
 
@@ -1215,12 +1224,14 @@ void std_be_build_stratum_request( char *req, struct work *work )
    bin2hex( ntimestr, (char*)(&ntime), sizeof(uint32_t) );
    bin2hex( noncestr, (char*)(&nonce), sizeof(uint32_t) );
    xnonce2str = abin2hex( work->xnonce2, work->xnonce2_len );
-   snprintf( req, JSON_BUF_LEN,
-        "{\"method\": \"mining.submit\", \"params\": [\"%s\", \"%s\", \"%s\", \"%s\", \"%s\"], \"id\":4}",
-         rpc_user, work->job_id, xnonce2str, ntimestr, noncestr );
+   snprintf( req, JSON_BUF_LEN, json_submit_req, rpc_user, work->job_id,
+             xnonce2str, ntimestr, noncestr );
    free( xnonce2str );
 }
 
+static const char *json_getwork_req = 
+  "{\"method\": \"getwork\", \"params\": [\"%s\"], \"id\":4}\r\n";
+
 bool std_le_submit_getwork_result( CURL *curl, struct work *work )
 {
    char req[JSON_BUF_LEN];
@@ -1237,8 +1248,7 @@ bool std_le_submit_getwork_result( CURL *curl, struct work *work )
       return false;
    }
    // build JSON-RPC request 
-   snprintf( req, JSON_BUF_LEN,
-     "{\"method\": \"getwork\", \"params\": [\"%s\"], \"id\":4}\r\n", gw_str );
+   snprintf( req, JSON_BUF_LEN, json_getwork_req, gw_str );
    free( gw_str );
    // issue JSON-RPC request 
    val = json_rpc_call( curl, rpc_url, rpc_userpass, req, NULL, 0 );
@@ -1271,8 +1281,7 @@ bool std_be_submit_getwork_result( CURL *curl, struct work *work )
       return false;
    }
    // build JSON-RPC request 
-   snprintf( req, JSON_BUF_LEN,
-     "{\"method\": \"getwork\", \"params\": [\"%s\"], \"id\":4}\r\n", gw_str );
+   snprintf( req, JSON_BUF_LEN, json_getwork_req, gw_str );
    free( gw_str );
    // issue JSON-RPC request 
    val = json_rpc_call( curl, rpc_url, rpc_userpass, req, NULL, 0 );
@@ -1327,25 +1336,33 @@ char* std_malloc_txs_request( struct work *work )
 
 static bool submit_upstream_work( CURL *curl, struct work *work )
 {
-   /* pass if the previous hash is not the current previous hash */
-   if ( !submit_old && memcmp( &work->data[1], &g_work.data[1], 32 ) )
-   {
-      if (opt_debug)
-         applog(LOG_DEBUG, "DEBUG: stale work detected, discarding");
-      return true;
-   }
-
+   work->stale = false;
+/*
+   // Submit anyway, discardring here messes up the stats
    if ( !have_stratum && allow_mininginfo )
    {
-      struct work wheight;
-      get_mininginfo( curl, &wheight );
-      if ( work->height && work->height <= net_blocks )
+      struct work mining_info;
+      get_mininginfo( curl, &mining_info );
+      if ( work->height < mining_info.height )
       {
-         if (opt_debug)
- 	        applog(LOG_WARNING, "block %u was already solved", work->height);
-	      return true;
+ 	      if ( !opt_quiet )
+            applog( LOG_WARNING, "Block %u already solved, current block %d",
+                                 work->height, mining_info.height );
+         work->stale = true;
+//         g_work_time = 0;
       }
    }
+*/
+
+   /* pass if the previous hash is not the current previous hash */
+   if ( !submit_old && ( work->data[ algo_gate.ntime_index ]
+                     != g_work.data[ algo_gate.ntime_index] ) )
+   {
+      if ( !opt_debug )
+         applog( LOG_WARNING, "Stale work detected, submitting anyway" );
+      work->stale = true;
+      g_work_time = 0;
+   }
 
    if ( have_stratum )
    {
@@ -1467,7 +1484,9 @@ start:
    else
       rc = work_decode( json_object_get( val, "result" ), work );
 
-   if ( opt_protocol && rc )
+   if ( rc ) 
+   {
+      if ( opt_protocol )
       {
          timeval_subtract( &diff, &tv_end, &tv_start );
          applog( LOG_DEBUG, "got new work in %.2f ms",
@@ -1477,6 +1496,70 @@ start:
       json_decref( val );
       // store work height in solo
       get_mininginfo(curl, work);
+
+      if ( work->height > last_block_height )
+      {
+         new_work_count++;
+         last_block_height = work->height;
+         applog( LOG_BLUE, "New block %d, net diff %.5g, target diff %.5g, ntime %08x",
+                 work->height, net_diff, work->targetdiff,
+                 bswap_32( work->data[ algo_gate.ntime_index ] ) );
+
+         if ( !opt_quiet && net_diff && net_hashrate )
+         {
+            double miner_hr = 0.;
+            pthread_mutex_lock( &stats_lock );
+
+            for ( int i = 0; i < opt_n_threads; i++ )
+               miner_hr += thr_hashrates[i];
+            global_hashrate = miner_hr;
+
+            pthread_mutex_unlock( &stats_lock );
+
+            if ( miner_hr )
+            {
+               double net_hr = net_hashrate;
+               char net_hr_units[4] = {0};
+               char miner_hr_units[4] = {0};
+               char net_ttf[32];
+               char miner_ttf[32];
+
+               sprintf_et( net_ttf, ( work->targetdiff * exp32 ) / net_hr );
+               sprintf_et( miner_ttf, ( work->targetdiff * exp32 ) / miner_hr );
+               scale_hash_for_display ( &miner_hr, miner_hr_units );
+               scale_hash_for_display ( &net_hr, net_hr_units );
+               applog2( LOG_INFO,
+                         "Miner TTF @ %.2f %sh/s %s, net TTF @ %.2f %sh/s %s",
+                                miner_hr, miner_hr_units, miner_ttf,
+                                net_hr, net_hr_units, net_ttf );
+            }
+         }
+      }  // work->height > last_block_height
+      else if ( memcmp( &work->data[1], &g_work.data[1], 32 ) )
+      {
+         new_work_count++;
+         applog( LOG_BLUE, "New work, ntime %08lx",
+                  bswap_32( work->data[ algo_gate.ntime_index ] ) );
+         if ( opt_debug )
+         {
+            uint32_t *old = g_work.data;
+            uint32_t *new = work->data;
+            printf("old: %08x %08x %08x %08x %08x %08x %08x %08x/n",
+              old[0],old[1],old[2],old[3],old[4],old[5],old[6],old[7]); 
+            printf("     %08x %08x %08x %08x %08x %08x %08x %08x/n",
+              old[8],old[9],old[10],old[11],old[12],old[13],old[14],old[15]); 
+            printf("     %08x %08x %08x %08x/n",
+                    old[16],old[17],old[18],old[19]); 
+            printf("new: %08x %08x %08x %08x %08x %08x %08x %08x/n",
+              new[0],new[1],new[2],new[3],new[4],new[5],new[6],new[7]);       
+            printf("     %08x %08x %08x %08x %08x %08x %08x %08x/n",
+              new[8],new[9],new[10],new[11],new[12],new[13],new[14],new[15]);  
+            printf("     %08x %08x %08x %08x/n",
+                    new[16],new[17],new[18],new[19]);                        
+         }
+      }
+   }  // rc
+
    return rc;
 }
 
@@ -1639,8 +1722,7 @@ static bool get_work(struct thr_info *thr, struct work *work)
 	return true;
 }
 
-bool submit_work( struct thr_info *thr,
-                         const struct work *work_in )
+bool submit_work( struct thr_info *thr, const struct work *work_in )
 {
 	struct workio_cmd *wc;
 
@@ -1664,27 +1746,44 @@ err_out:
 	return false;
 }
 
+/*
 // __float128?
 // Convert little endian 256 bit (38 decimal digits) unsigned integer to
 // double precision floating point with 15 decimal digits precision.
-// returns u * ( 2**256 ) 
 static inline double u256_to_double( const uint64_t *u )
 {
-   const double exp64 = 4294967296.0 * 4294967296.0;  // 2**64
    return ( ( u[3] * exp64 + u[2] ) * exp64 + u[1] ) * exp64 + u[0];
 }
+*/
 
+static void update_submit_stats( struct work *work, const void *hash )
+{
+   work->sharediff = hash ? exp32 / ( (uint64_t*)hash )[3] : 0.;
+
+   pthread_mutex_lock( &stats_lock );
+
+   submitted_share_count++;
+   share_stats[ s_put_ptr ].share_count = submitted_share_count;
+   gettimeofday( &share_stats[ s_put_ptr ].submit_time, NULL );
+   share_stats[ s_put_ptr ].share_diff = work->sharediff;
+   share_stats[ s_put_ptr ].net_diff = net_diff;
+   share_stats[ s_put_ptr ].stratum_diff = stratum_diff;
+   share_stats[ s_put_ptr ].target_diff = work->targetdiff;
+   if ( have_stratum )
+      strncpy( share_stats[ s_put_ptr ].job_id, work->job_id, 30 );
+   s_put_ptr = stats_ptr_incr( s_put_ptr );
+
+   pthread_mutex_unlock( &stats_lock );
+}
+
+/*
+//deprecated
 void work_set_target_ratio( struct work* work, const void *hash )
 {
-   if ( likely( hash ) )
-   {
-      double dhash = u256_to_double( (const uint64_t*)hash );
-      if ( likely( dhash > 0. ) )
-         work->sharediff = work->targetdiff *
-             u256_to_double( (const uint64_t*)( work->target ) ) / dhash;
-   }
-   else
-      work->sharediff = 0.;
+   submitted_share_count++;
+   work->sharediff = work->targetdiff * (double)( ((uint64_t*)hash)[3] )
+                                / (double)( ((uint64_t*)work->target)[3] );
+//   work->sharediff = likely( hash ) ? target_to_diff( (uint32_t*)hash ) : 0.;
 
    // collect some share stats
    // Frequent share submission combined with high latency can caused
@@ -1698,25 +1797,33 @@ void work_set_target_ratio( struct work* work, const void *hash )
    share_stats[ s_put_ptr ].net_diff = net_diff;
    share_stats[ s_put_ptr ].stratum_diff = stratum_diff;
    share_stats[ s_put_ptr ].target_diff = work->targetdiff;
-   ( (uint64_t*)share_stats[ s_put_ptr ].job_id )[3] = 0;
+   if ( have_stratum )
       strncpy( share_stats[ s_put_ptr ].job_id, work->job_id, 30 );
    s_put_ptr = stats_ptr_incr( s_put_ptr );
 
    pthread_mutex_unlock( &stats_lock );
 }
+*/
 
 bool submit_solution( struct work *work, const void *hash,
                       struct thr_info *thr )
 {
   if ( likely( submit_work( thr, work ) ) )
   {
-     submitted_share_count++;
-     work_set_target_ratio( work, hash );
-     if ( !opt_quiet )
-        applog( LOG_NOTICE, "%d submitted by thread %d, job %s",
-            submitted_share_count, thr->id, work->job_id );
+     update_submit_stats( work, hash );
 
-if ( lowdiff_debug )
+     if ( !opt_quiet )
+     {
+            if ( have_stratum )
+              applog( LOG_NOTICE, "%d Submit diff %.5g, block %d, job %s",
+                      submitted_share_count, work->sharediff, work->height,
+                      work->job_id );
+            else
+              applog( LOG_NOTICE, "%d Submit diff %.5g, block %d",
+                      submitted_share_count, work->sharediff, work->height );
+     }
+
+     if ( unlikely( lowdiff_debug ) )
      {
         uint32_t* h = (uint32_t*)hash;
         uint32_t* t = (uint32_t*)work->target;
@@ -1728,21 +1835,28 @@ if ( lowdiff_debug )
      return true;
   }
   else
-     applog( LOG_WARNING, "%d failed to submit share.",
-             submitted_share_count );
+     applog( LOG_WARNING, "%d failed to submit share", submitted_share_count );
   return false;
 }
 
+// deprecated, use submit_solution
 bool submit_lane_solution( struct work *work, const void *hash,
                            struct thr_info *thr, const int lane )
 {
   if ( likely( submit_work( thr, work ) ) )
   {
-     submitted_share_count++;
-     work_set_target_ratio( work, hash );
+     update_submit_stats( work, hash );
+
      if ( !opt_quiet )
-        applog( LOG_NOTICE, "%d submitted by thread %d, lane %d, job %s",
-            submitted_share_count, thr->id, lane, work->job_id );
+     {
+            if ( have_stratum )
+              applog( LOG_NOTICE, "%d Submit diff %.5g, block %d, job %s",
+                      submitted_share_count, work->sharediff, work->height,
+                      work->job_id );
+            else
+              applog( LOG_NOTICE, "%d Submit diff %.5g, block %d",
+                      submitted_share_count, work->sharediff, work->height );
+     }
 
      if ( lowdiff_debug )
      {
@@ -1753,13 +1867,11 @@ if ( lowdiff_debug )
          applog(LOG_INFO,"Targ[7:0]: %08x %08x %08x %08x %08x %08x %08x %08x",
                                      t[7],t[6],t[5],t[4],t[3],t[2],t[1],t[0]);
     }
-
-
     return true;
   }
   else
-     applog( LOG_WARNING, "%d failed to submit share.",
-          submitted_share_count );
+     applog( LOG_WARNING, "%d failed to submit share", submitted_share_count );
+
   return false;
 }
 
@@ -1853,12 +1965,15 @@ void std_get_new_work( struct work* work, struct work* g_work, int thr_id,
                      uint32_t *end_nonce_ptr )
 {
    uint32_t *nonceptr = work->data + algo_gate.nonce_index;
+   bool force_new_work = false; 
 
-   bool force_new_work = work->job_id ? strtoul(   work->job_id, NULL, 16 ) !=
-                                        strtoul( g_work->job_id, NULL, 16 )
-                                      : true;
+   if ( have_stratum ) 
+      force_new_work = work->job_id ?    strtoul(   work->job_id, NULL, 16 )
+                                      != strtoul( g_work->job_id, NULL, 16 )
+                                     : false;
 
-   if ( force_new_work || *nonceptr >= *end_nonce_ptr )
+   if ( force_new_work || ( *nonceptr >= *end_nonce_ptr )
+     || memcmp( work->data, g_work->data, algo_gate.work_cmp_size ) )
    {
      work_free( work );
      work_copy( work, g_work );
@@ -2006,12 +2121,13 @@ static void *miner_thread( void *userdata )
           }
           else
           {
-             int min_scantime = have_longpoll ? LP_SCANTIME : opt_scantime;
+             int scantime = have_longpoll ? LP_SCANTIME : opt_scantime;
 	          pthread_mutex_lock( &g_work_lock );
 
-             if ( time(NULL) - g_work_time >= min_scantime
+             if ( time(NULL) - g_work_time >= scantime
                   || *nonceptr >= end_nonce )
              {
+                report_summary_log( false );
                 if ( unlikely( !get_work( mythr, &g_work ) ) )
                 {
 		             applog( LOG_ERR, "work retrieval failed, exiting "
@@ -2020,6 +2136,7 @@ static void *miner_thread( void *userdata )
 		             goto out;
 	             }
                 g_work_time = time(NULL);
+                restart_threads();
              }
              algo_gate.get_new_work( &work, &g_work, thr_id, &end_nonce );
 
@@ -2036,12 +2153,28 @@ static void *miner_thread( void *userdata )
           sleep(5);
 	       continue;
        }
-       // adjust max_nonce to meet target scan time
+
+// LP_SCANTIME overrides opt_scantime option, is this right?
+
+       // adjust max_nonce to meet target scan time. Startum and longpoll
+       // can go longer because they can rely on restart_threads to signal
+       // an early abort. get_work on the other hand can't rely on
+       // restart_threads so need a much shorter scantime
+       if ( have_stratum )
+          max64 = 60 * thr_hashrates[thr_id];
+       else if ( have_longpoll )
+          max64 = LP_SCANTIME * thr_hashrates[thr_id];
+       else  // getwork inline
+          max64 = opt_scantime * thr_hashrates[thr_id];   
+
+/*       
        if ( have_stratum )
           max64 = LP_SCANTIME;
        else
           max64 = g_work_time + ( have_longpoll ? LP_SCANTIME : opt_scantime )
 	                      - time(NULL);
+*/
+
        // time limit
        if ( unlikely( opt_time_limit && firstwork_time ) )
        {
@@ -2068,17 +2201,20 @@ static void *miner_thread( void *userdata )
           }
           if ( remain < max64 ) max64 = remain;
        }
-       // Select nonce range for approx 1 min duration based
-       // on hashrate, initial value arbitrarilly set to 1000 just to get
+
+       // Select nonce range based on max64, the estimated number of hashes
+       // to meet the desired scan time.
+       // Initial value arbitrarilly set to 1000 just to get
        // a sample hashrate for the next time.
        uint32_t work_nonce = *nonceptr;
-       max64 = 60 * thr_hashrates[thr_id];
+//       max64 = 60 * thr_hashrates[thr_id];
        if ( max64 <= 0)
           max64 = 1000;
        if ( work_nonce + max64 > end_nonce )
           max_nonce = end_nonce;
        else
           max_nonce = work_nonce + (uint32_t)max64;
+
        // init time
        if ( firstwork_time == 0 )
           firstwork_time = time(NULL);
@@ -2148,7 +2284,7 @@ static void *miner_thread( void *userdata )
              }
              else
                 sprintf( tempstr, "%d C", temp );
-             applog( LOG_INFO,"CPU temp: curr %s (max %d), Freq: %.3f/%.3f GHz",
+             applog( LOG_NOTICE,"CPU temp: curr %s (max %d), Freq: %.3f/%.3f GHz",
                      tempstr, hi_temp, (float)lo_freq / 1e6, (float)hi_freq/ 1e6 );
              if ( temp > hi_temp ) hi_temp = temp;
           }
@@ -2209,6 +2345,8 @@ void restart_threads(void)
 {
 	for ( int i = 0; i < opt_n_threads; i++)
 		work_restart[i].restart = 1;
+   if ( opt_debug )
+      applog( LOG_INFO, "Threads restarted for new work."); 
 }
 
 json_t *std_longpoll_rpc_call( CURL *curl, int *err, char* lp_url )
@@ -2292,6 +2430,10 @@ start:
       soval = json_object_get(res, "submitold");
       submit_old = soval ? json_is_true(soval) : false;
 	   pthread_mutex_lock(&g_work_lock);
+
+// This code has been here for a long time even though job_id isn't used.
+// This needs to be changed eventually to test the block height properly
+// using g_work.block_height .     
       start_job_id = g_work.job_id ? strdup(g_work.job_id) : NULL;
 	   if (have_gbt)
 	      rc = gbt_work_decode(res, &g_work);
@@ -2299,6 +2441,7 @@ start:
 	      rc = work_decode(res, &g_work);
 	   if (rc)
       {
+// purge job id from solo mining
         bool newblock = g_work.job_id && strcmp(start_job_id, g_work.job_id);
 	     newblock |= (start_diff != net_diff); // the best is the height but... longpoll...
         if (newblock)
@@ -2459,8 +2602,10 @@ void std_stratum_gen_work( struct stratum_ctx *sctx, struct work *g_work )
    algo_gate.build_extraheader( g_work, sctx );
    net_diff = algo_gate.calc_network_diff( g_work );
    algo_gate.set_work_data_endian( g_work );
-   work_set_target( g_work, sctx->job.diff
-                                  / ( opt_target_factor * opt_diff_factor ) );
+   g_work->height = sctx->block_height;
+   g_work->targetdiff = sctx->job.diff
+                                   / ( opt_target_factor * opt_diff_factor );
+   diff_to_target( g_work->target, g_work->targetdiff );
 
    pthread_mutex_unlock( &sctx->work_lock );
 
@@ -2484,13 +2629,16 @@ void std_stratum_gen_work( struct stratum_ctx *sctx, struct work *g_work )
    pthread_mutex_unlock( &stats_lock );
 
    if ( stratum_diff != sctx->job.diff )
-      applog( LOG_BLUE, "New stratum diff %g, block %d, job %s",
-                        sctx->job.diff, sctx->block_height, g_work->job_id );
+      applog( LOG_BLUE, "New diff %g, block %d, job %s, ntime %08x",
+                        sctx->job.diff, sctx->block_height, g_work->job_id,
+                        bswap_32( g_work->data[ algo_gate.ntime_index ] ) );
    else if ( last_block_height != sctx->block_height )
-      applog( LOG_BLUE, "New block %d, job %s",
-                         sctx->block_height, g_work->job_id );
+      applog( LOG_BLUE, "New block %d, job %si, ntime %08x",
+                        sctx->block_height, g_work->job_id,
+                        bswap_32( g_work->data[ algo_gate.ntime_index ] ) );
    else if ( g_work->job_id ) 
-      applog( LOG_BLUE,"New job %s", g_work->job_id );
+      applog( LOG_BLUE,"New job %s, ntime %08x", g_work->job_id,
+                        bswap_32( g_work->data[ algo_gate.ntime_index ] ) );
 
    // Update data and calculate new estimates.
    if ( ( stratum_diff != sctx->job.diff )
@@ -2506,9 +2654,8 @@ void std_stratum_gen_work( struct stratum_ctx *sctx, struct work *g_work )
 
       if ( !opt_quiet )
       {
-         applog2( LOG_INFO, "%s: %s", algo_names[opt_algo], short_url );
          applog2( LOG_INFO, "Diff: Net %.5g, Stratum %.5g, Target %.5g",
-                            net_diff, stratum_diff, last_targetdiff );
+                            net_diff, stratum_diff, g_work->targetdiff );
 
          if ( likely( hr > 0. ) )
          {
@@ -2516,8 +2663,8 @@ void std_stratum_gen_work( struct stratum_ctx *sctx, struct work *g_work )
             char block_ttf[32];
             char share_ttf[32];
 
-            sprintf_et( block_ttf, net_diff * diff_to_hash / hr );
-            sprintf_et( share_ttf, last_targetdiff * diff_to_hash / hr );
+            sprintf_et( block_ttf, ( net_diff * exp32 ) /  hr );
+            sprintf_et( share_ttf, g_work->targetdiff * exp32 / hr );
             scale_hash_for_display ( &hr, hr_units );
             applog2( LOG_INFO, "TTF @ %.2f %sh/s: block %s, share %s",
                                hr, hr_units, block_ttf, share_ttf );
@@ -2532,7 +2679,7 @@ void std_stratum_gen_work( struct stratum_ctx *sctx, struct work *g_work )
                     : et.tv_sec / ( last_block_height - session_first_block );
                if ( net_diff && net_ttf )
                {
-                  double net_hr = net_diff * diff_to_hash / net_ttf;
+                  double net_hr = net_diff * exp32 / net_ttf;
                   char net_ttf_str[32];
                   char net_hr_units[4] = {0};
 
@@ -3594,9 +3741,6 @@ int main(int argc, char *argv[])
 */
    }
 
-   applog( LOG_INFO, "Extranonce subscribe: %s",
-                           opt_extranonce ? "YES" : "NO" );
-
 #ifdef HAVE_SYSLOG_H
 	if (use_syslog)
 		openlog("cpuminer", LOG_PID, LOG_USER);
@@ -3634,6 +3778,9 @@ int main(int argc, char *argv[])
 	/* ESET-NOD32 Detects these 2 thread_create... */
 	if (want_longpoll && !have_stratum)
    {
+      if ( opt_debug )
+         applog(LOG_INFO,"Creating long poll thread");
+
       /* init longpoll thread info */
 		longpoll_thr_id = opt_n_threads + 1;
 		thr = &thr_info[longpoll_thr_id];
@@ -3650,6 +3797,9 @@ int main(int argc, char *argv[])
 	}
 	if (want_stratum)
    {
+      if ( opt_debug )
+         applog(LOG_INFO,"Creating stratum thread");
+
       /* init stratum thread info */
 		stratum_thr_id = opt_n_threads + 2;
 		thr = &thr_info[stratum_thr_id];
@@ -3670,6 +3820,9 @@ int main(int argc, char *argv[])
 
 	if ( opt_api_enabled )
    {
+      if ( opt_debug )
+         applog(LOG_INFO,"Creating API thread");
+
       /* api thread */
 		api_thr_id = opt_n_threads + 3;
 		thr = &thr_info[api_thr_id];
@@ -3680,7 +3833,7 @@ int main(int argc, char *argv[])
 		err = thread_create( thr, api_thread );
 		if ( err )
       {
-			applog( LOG_ERR, "api thread create failed" );
+			applog( LOG_ERR, "API thread create failed" );
 			return 1;
 		}
       if ( !opt_quiet )

miner.h

@@ -312,6 +312,19 @@ int    varint_encode( unsigned char *p, uint64_t n );
 size_t address_to_script( unsigned char *out, size_t outsz, const char *addr );
 int    timeval_subtract( struct timeval *result, struct timeval *x,
                            struct timeval *y);
+
+// Bitcoin formula for converting difficulty to an equivalent
+// number of hashes.
+//
+//     https://en.bitcoin.it/wiki/Difficulty
+//
+//     hash = diff * 2**32
+//
+// diff_to_hash = 2**32 = 0x100000000 = 4294967296 = exp32;
+
+const double exp32;  // 2**32
+const double exp64;  // 2**64
+
 bool   fulltest( const uint32_t *hash, const uint32_t *target );
 bool   valid_hash( const void*, const void* );
 
@@ -334,10 +347,11 @@ struct thr_info {
 
 bool submit_solution( struct work *work, const void *hash,
                       struct thr_info *thr );
+
+// deprecated
 bool submit_lane_solution( struct work *work, const void *hash,
                            struct thr_info *thr, const int lane );
 
-
 bool submit_work( struct thr_info *thr, const struct work *work_in );
 
 
@@ -378,6 +392,7 @@ struct work {
 	size_t xnonce2_len;
 	unsigned char *xnonce2;
    bool sapling;
+   bool stale;
 
    // x16rt
    uint32_t merkleroothash[8];
@@ -854,7 +869,7 @@ Options:\n\
                           zr5           Ziftr\n\
   -N, --param-n         N parameter for scrypt based algos\n\
   -R, --patam-r         R parameter for scrypt based algos\n\
-  -K, --param-key       Key parameter for algos that use it\n\
+  -K, --param-key       Key (pers) parameter for algos that use it\n\
   -o, --url=URL         URL of mining server\n\
   -O, --userpass=U:P    username:password pair for mining server\n\
   -u, --user=USERNAME   username for mining server\n\

util.c

@@ -983,6 +983,7 @@ int timeval_subtract(struct timeval *result, struct timeval *x,
 	return x->tv_sec < y->tv_sec;
 }
 
+// deprecated, use test_hash_and_submit
 // Use this when deinterleaved
 // do 64 bit test 4 iterations
 inline bool valid_hash( const void *hash, const void *target )
@@ -999,6 +1000,7 @@ inline bool valid_hash( const void *hash, const void *target )
    return true;
 }
 
+// deprecated, use test_hash_and_submit
 bool fulltest( const uint32_t *hash, const uint32_t *target )
 {
 	int i;
@@ -1044,32 +1046,42 @@ void diff_to_target(uint32_t *target, double diff)
    uint64_t m;
    int k;
 
-   const double exp64 = (double)0xffffffffffffffff + 1.;   
-   for ( k = 3; k > 0 && diff > 1.0; k-- )
-      diff /= exp64;
+   for (k = 6; k > 0 && diff > 1.0; k--)
+      diff /= exp32;
 
-//   for (k = 6; k > 0 && diff > 1.0; k--)
 //      diff /= 4294967296.0;
-	m = (uint64_t)( 0xffff0000 / diff );
-	if unlikely( m == 0 && k == 3 )
+
+//   m = (uint64_t)(4294901760.0 / diff);
+
+   m = (uint64_t)(exp32 / diff);
+
+   if (m == 0 && k == 6)
       memset(target, 0xff, 32);
-	else
-   {
+   else {
       memset(target, 0, 32);
-      ((uint64_t*)target)[k] = m;
-//		target[k] = (uint32_t)m;
-//		target[k + 1] = (uint32_t)(m >> 32);
+      target[k] = (uint32_t)m;
+      target[k + 1] = (uint32_t)(m >> 32);
    }
 }
 
-// Only used by stratum pools
+
+// deprecated
 void work_set_target(struct work* work, double diff)
 {
 	diff_to_target( work->target, diff );
 	work->targetdiff = diff;
 }
 
-// Only used by longpoll pools
+double target_to_diff( uint32_t* target )
+{
+   uint64_t *targ = (uint64_t*)target;
+   return target ? 1. / ( ( (double)targ[3] / exp32 )
+                        + ( (double)targ[2]         )
+                        + ( (double)targ[1] * exp32 )
+                        + ( (double)targ[0] * exp64 ) )
+                  : 0.;       
+}
+/*
 double target_to_diff(uint32_t* target)
 {
 	uchar* tgt = (uchar*) target;
@@ -1088,6 +1100,7 @@ double target_to_diff(uint32_t* target)
 	else
 		return (double)0x0000ffff00000000/m;
 }
+*/
 
 #ifdef WIN32
 #define socket_blocks() (WSAGetLastError() == WSAEWOULDBLOCK)
@@ -1555,18 +1568,27 @@ bool stratum_authorize(struct stratum_ctx *sctx, const char *user, const char *p
 	if ( !stratum_send_line( sctx, s ) )
 		goto out;
 
-	if (!socket_full(sctx->sock, 3)) {
-         applog(LOG_WARNING, "stratum extranonce subscribe timed out");
+	if ( !socket_full( sctx->sock, 3 ) )
+   {
+      applog( LOG_WARNING, "Extranonce disabled, subscribe timed out" );
+		opt_extranonce = false;
       goto out;
 	}
+   if ( !opt_quiet )
+      applog( LOG_INFO, "Extranonce subscription enabled" );
 
 	sret = stratum_recv_line( sctx );
-	if (sret) {
+	if ( sret )
+   {
 		json_t *extra = JSON_LOADS( sret, &err );
-		if (!extra) {
+		if ( !extra )
+      {
 			applog(LOG_WARNING, "JSON decode failed(%d): %s", err.line, err.text);
-		} else {
-			if (json_integer_value(json_object_get(extra, "id")) != 3) {
+		}
+      else
+      {
+			if ( json_integer_value(json_object_get( extra, "id" ) ) != 3 )
+         {
 				// we receive a standard method if extranonce is ignored
 				if ( !stratum_handle_method( sctx, sret ) )
 					applog( LOG_WARNING, "Stratum answer id is not correct!" );