Release v1.4 with blake (NEOS)

Blake256: squashed commit...

Squashed commit of the following:

commit c370208bc92ef16557f66e5391faf2b1ad47726f
Author: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date:   Wed Sep 3 13:53:01 2014 +0200

    hashlog: prepare store of scanned range

commit e2cf49a5e956f03deafd266d1a0dd087a2041c99
Author: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date:   Wed Sep 3 12:54:13 2014 +0200

    stratum: store server time offset in context

commit 1a4391d7ff
Author: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date:   Tue Sep 2 12:40:52 2014 +0200

    hashlog: prevent double computing on jobs already done

commit 049e577301
Author: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date:   Wed Sep 3 09:49:14 2014 +0200

    tmp blake log

commit 43d3e93e1a
Author: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date:   Wed Sep 3 09:29:51 2014 +0200

    blake: set a max throughput

commit 7e595a36ea
Author: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date:   Tue Sep 2 21:13:37 2014 +0200

    blake: cleanup, remove d_hash buf, not in a chain

    host: only bencode if gpu hash was found

commit de80c7e9d1
Author: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date:   Tue Sep 2 12:40:44 2014 +0200

    blake: remove unused parameter and fix index in d_hash

    that reduce the speed to 92MH/s but the next commit
    give us 30 more

    so, todo: merge the whole checkhash proc in gpu_hash
              and remove this d_hash buffer...

commit 2d42ae6de5
Author: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date:   Tue Sep 2 05:09:31 2014 +0200

    stratum: handle a small cache of submitted jobs

    Prevent to send duplicated shares on some pools like hashharder..

    This cache keeps submitted job/nounces of the last 15 minutes

    so, remove exit on repeated duplicate shares,
        the submitted cache now handles this problem.

    Signed-off-by: Tanguy Pruvot <tanguy.pruvot@gmail.com>

commit 1b8c3c12fa
Author: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date:   Tue Sep 2 03:38:57 2014 +0200

    debug: a new boolean to log or not json rpc data

commit 1f99aae0ff
Author: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date:   Mon Sep 1 18:49:23 2014 +0200

    exit on repeated duplicate shares (to enhance)

    create a new function proper_exit() to do common stuff on exit...

commit 530732458a
Author: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date:   Mon Sep 1 12:22:51 2014 +0200

    blake: use a constant for threads, reduce mallocated d_hash size

    and clean a bit more...

commit 0aeac878ef
Author: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date:   Mon Sep 1 06:12:55 2014 +0200

    blake: tune up and cleanup, ~100 MH/s on a normal 750Ti

    tested on linux and windows (x86 binary)...

    but there is a high number of duplicated shares... weird

commit 4a52d0553b
Author: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date:   Mon Sep 1 10:22:32 2014 +0200

    debug: show json methods, hide hash/target if ok

commit 1fb9becc1f
Author: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date:   Mon Sep 1 08:44:19 2014 +0200

    cpu-miner: sort algos by name, show reject reason

commit bfe96c49b0
Author: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date:   Mon Aug 25 11:21:06 2014 +0200

    release 1.4, update README...

commit c17d11e377
Author: Tanguy Pruvot <tanguy.pruvot@gmail.com>
Date:   Sun Aug 31 08:57:48 2014 +0200

    add "blake" 256, 14 rounds (for NEOS blake, not BlakeCoin)

    also remove "missing" file, its old and not compatible with ubuntu 14.04

to test on windows

blake: clean and optimize

Release v1.4 with blake (NEOS)

Makefile.am

@@ -17,6 +17,7 @@ ccminer_SOURCES		= elist.h miner.h compat.h \
 			  compat/inttypes.h compat/stdbool.h compat/unistd.h \
 			  compat/sys/time.h compat/getopt/getopt.h \
 			  cpu-miner.c util.c hefty1.c scrypt.c \
+			  hashlog.cpp \
 			  heavy/heavy.cu \
 			  heavy/cuda_blake512.cu heavy/cuda_blake512.h \
 			  heavy/cuda_combine.cu heavy/cuda_combine.h \
@@ -32,7 +33,7 @@ ccminer_SOURCES		= elist.h miner.h compat.h \
 			  quark/cuda_jh512.cu quark/cuda_quark_blake512.cu quark/cuda_quark_groestl512.cu quark/cuda_skein512.cu \
 			  quark/cuda_bmw512.cu quark/cuda_quark_keccak512.cu quark/quarkcoin.cu quark/animecoin.cu \
 			  quark/cuda_quark_compactionTest.cu \
-			  cuda_nist5.cu \
+			  cuda_nist5.cu blake32.cu \
 			  sph/bmw.c sph/blake.c sph/groestl.c sph/jh.c sph/keccak.c sph/skein.c \
 			  sph/cubehash.c sph/echo.c sph/luffa.c sph/sha2.c sph/shavite.c sph/simd.c \
 			  sph/hamsi.c sph/hamsi_helper.c sph/sph_hamsi.h \
@@ -43,6 +44,7 @@ ccminer_SOURCES		= elist.h miner.h compat.h \
 			  x15/x14.cu x15/x15.cu x15/cuda_x14_shabal512.cu x15/cuda_x15_whirlpool.cu x15/whirlpool.cu \
 			  x17/x17.cu x17/cuda_x17_haval512.cu x17/cuda_x17_sha512.cu
 
+
 ccminer_LDFLAGS		= $(PTHREAD_FLAGS) @CUDA_LDFLAGS@
 ccminer_LDADD		= @LIBCURL@ @JANSSON_LIBS@ @PTHREAD_LIBS@ @WS2_LIBS@ @CUDA_LIBS@ @OPENMP_CFLAGS@ @LIBS@
 ccminer_CPPFLAGS	= -msse2 @LIBCURL_CPPFLAGS@ @OPENMP_CFLAGS@ $(PTHREAD_FLAGS) -fno-strict-aliasing $(JANSSON_INCLUDES) -DSCRYPT_KECCAK512 -DSCRYPT_CHACHA -DSCRYPT_CHOOSE_COMPILETIME
@@ -51,30 +53,33 @@ nvcc_ARCH  = -gencode=arch=compute_50,code=\"sm_50,compute_50\"
 #nvcc_ARCH += -gencode=arch=compute_35,code=\"sm_35,compute_35\"
 #nvcc_ARCH += -gencode=arch=compute_30,code=\"sm_30,compute_30\"
 
-nvcc_FLAGS = $(nvcc_ARCH) -I . --ptxas-options=-v --use_fast_math
+nvcc_FLAGS = $(nvcc_ARCH) -I . @CUDA_CFLAGS@
 nvcc_FLAGS += $(JANSSON_INCLUDES)
 
 # we're now targeting all major compute architectures within one binary.
 .cu.o:
-	$(NVCC) $(nvcc_FLAGS) @CFLAGS@ --maxrregcount=128 -o $@ -c $<
+	$(NVCC) $(nvcc_FLAGS) --maxrregcount=128 -o $@ -c $<
+
+blake32.o: blake32.cu
+	$(NVCC) $(nvcc_FLAGS) --maxrregcount=64 -o $@ -c $<
 
 # Luffa and Echo are faster with 80 registers than 128
 x11/cuda_x11_luffa512.o: x11/cuda_x11_luffa512.cu
-	$(NVCC) $(nvcc_FLAGS) @CFLAGS@ --maxrregcount=80 -o $@ -c $<
+	$(NVCC) $(nvcc_FLAGS) --maxrregcount=80 -o $@ -c $<
 
 x11/cuda_x11_echo.o: x11/cuda_x11_echo.cu
-	$(NVCC) $(nvcc_FLAGS) @CFLAGS@ --maxrregcount=80 -o $@ -c $<
+	$(NVCC) $(nvcc_FLAGS) --maxrregcount=80 -o $@ -c $<
 
 # Shavite compiles faster with 128 regs
 x11/cuda_x11_shavite512.o: x11/cuda_x11_shavite512.cu
-	$(NVCC) $(nvcc_FLAGS) -I cudpp-2.1/include @CFLAGS@ --maxrregcount=128 -o $@ -c $<
+	$(NVCC) $(nvcc_FLAGS) -I cudpp-2.1/include --maxrregcount=128 -o $@ -c $<
 
 x17/cuda_x17_sha512.o: x17/cuda_x17_sha512.cu
-	$(NVCC) $(nvcc_FLAGS) -O2 --maxrregcount=80 -o $@ -c $<
+	$(NVCC) $(nvcc_FLAGS) --maxrregcount=80 -o $@ -c $<
 
 # ABI requiring code modules
 quark/cuda_quark_compactionTest.o: quark/cuda_quark_compactionTest.cu
-	$(NVCC) $(nvcc_FLAGS) -I cudpp-2.1/include @CFLAGS@ -Xptxas "-abi=yes -v" --maxrregcount=80 -o $@ -c $<
+	$(NVCC) $(nvcc_FLAGS) -I cudpp-2.1/include --maxrregcount=80 -o $@ -c $<
 
 JHA/cuda_jha_compactionTest.o: JHA/cuda_jha_compactionTest.cu
-	$(NVCC) $(nvcc_FLAGS) -I cudpp-2.1/include @CFLAGS@ -Xptxas "-abi=yes -v" --maxrregcount=80 -o $@ -c $<
+	$(NVCC) $(nvcc_FLAGS) -I cudpp-2.1/include --maxrregcount=80 -o $@ -c $<

README.md

@@ -3,5 +3,32 @@ ccminer
 
 Christian Buchner's & Christian H.'s CUDA miner project
 
-Fork by tpruvot@github with X14 support
+Fork by tpruvot@github with X14,X15,X17,WHIRL and Blake256 support
+
    BTC donation address: 1AJdfCpLWPNoAMDfHF1wD5y8VgKSSTHxPo
+   [![tip for next commit](https://tip4commit.com/projects/927.svg)](https://tip4commit.com/github/tpruvot/ccminer)
+
+A big part of my recent additions were wrote by [djm34](https://github.com/djm34),
+You can also donate some beers (or redbulls)
+
+This variant was tested and built on Linux (ubuntu server 14.04)
+and VStudio 2013.
+
+Note that the x86 releases are faster than x64 ones on Windows.
+
+About source code dependencies
+------------------------------
+
+This project requires some libraries to be built :
+
+- OpenSSL
+
+- Curl
+
+- pthreads
+
+You can download prebuilt .lib and dll on the [bitcointalk forum thread](https://bitcointalk.org/?topic=167229.0)
+
+
+There is also a [Tutorial for windows](http://cudamining.co.uk/url/tutorials/id/3) on [CudaMining](http://cudamining.co.uk) website.
+

README.txt

@@ -63,11 +63,14 @@ its command line interface and options.
                           jackpot     use to mine Jackpotcoin
                           quark       use to mine Quarkcoin
                           anime       use to mine Animecoin
+                          blake       use to mine NEOS (Blake 256)
                           nist5       use to mine TalkCoin
                           fresh       use to mine Freshcoin
+                          whirl       use to mine Whirlcoin
                           x11         use to mine DarkCoin
                           x14         use to mine X14Coin
                           x15         use to mine Halcyon
+                          x17         use to mine X17
 
   -d, --devices         gives a comma separated list of CUDA device IDs
                         to operate on. Device IDs start counting from 0!
@@ -98,6 +101,7 @@ its command line interface and options.
       --benchmark       run in offline benchmark mode
       --cputest         debug hashes from cpu algorithms
   -c, --config=FILE     load a JSON-format configuration file
+  -C, --color           display colored output in a linux Terminal
   -V, --version         display version information and exit
   -h, --help            display this help text and exit
 
@@ -148,6 +152,14 @@ features.
 
 >>> RELEASE HISTORY <<<
 
+  Sep.  1st 2014  add X17, optimized x15 and whirl
+                  add blake (256 variant)
+                  color support on Windows,
+                  remove some dll dependencies (pthreads, msvcp)
+
+  Aug. 18th 2014  add X14, X15, Whirl, and Fresh algos,
+                  also add colors and nvprof cmd line support
+
   June 15th 2014  add X13 and Diamond Groestl support.
                   Thanks to tsiv and to Bombadil for the contributions!
 
@@ -214,6 +226,9 @@ Notable contributors to this application are:
 
 Christian Buchner, Christian H. (Germany): CUDA implementation 
 
+Tanguy Pruvot : CUDA, blake, general code cleanup, tuneup for linux (Makefiles)
+                and some vstudio 2013 stuff...
+
 and also many thanks to anyone else who contributed to the original
 cpuminer application (Jeff Garzik, pooler), it's original HVC-fork
 and the HVC-fork available at hvc.1gh.com

blake32.cu

@@ -0,0 +1,351 @@
+/**
+ * Blake-256 Cuda Kernel (Tested on SM 5.0)
+ *
+ * Tanguy Pruvot - Aug. 2014
+ */
+
+#include "miner.h"
+
+extern "C" {
+#include "sph/sph_blake.h"
+#include <stdint.h>
+#include <memory.h>
+}
+
+/* threads per block */
+#define TPB 128
+
+/* hash by cpu with blake 256 */
+extern "C" void blake32hash(void *output, const void *input)
+{
+	unsigned char hash[64];
+	sph_blake256_context ctx;
+	sph_blake256_init(&ctx);
+	sph_blake256(&ctx, input, 80);
+	sph_blake256_close(&ctx, hash);
+	memcpy(output, hash, 32);
+}
+
+#include "cuda_helper.h"
+
+// in cpu-miner.c
+extern bool opt_n_threads;
+extern bool opt_benchmark;
+extern int device_map[8];
+
+extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);
+
+__constant__
+static uint32_t __align__(32) c_PaddedMessage80[32]; // padded message (80 bytes + padding)
+
+__constant__
+static uint32_t __align__(32) c_Target[8];
+
+#define MAXU 0xffffffffU
+
+static uint32_t *d_resNounce[8];
+static uint32_t *h_resNounce[8];
+
+__constant__
+#ifdef WIN32
+/* what the fuck ! */
+static uint8_t c_sigma[16][16];
+const uint8_t host_sigma[16][16] =
+#else
+/* prefer uint32_t to prevent size conversions = speed +5/10 % */
+static uint32_t __align__(32) c_sigma[16][16];
+const uint32_t host_sigma[16][16]
+#endif
+= {
+	{ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
+	{14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
+	{11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
+	{ 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
+	{ 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 },
+	{ 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 },
+	{12, 5, 1, 15, 14, 13, 4, 10, 0, 7, 6, 3, 9, 2, 8, 11 },
+	{13, 11, 7, 14, 12, 1, 3, 9, 5, 0, 15, 4, 8, 6, 2, 10 },
+	{ 6, 15, 14, 9, 11, 3, 0, 8, 12, 2, 13, 7, 1, 4, 10, 5 },
+	{10, 2, 8, 4, 7, 6, 1, 5, 15, 11, 9, 14, 3, 12, 13 , 0 },
+	{ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
+	{14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
+	{11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
+	{ 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
+	{ 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 },
+	{ 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 }
+};
+
+__device__ __constant__
+static const uint32_t __align__(32) c_IV256[8] = {
+	SPH_C32(0x6A09E667), SPH_C32(0xBB67AE85),
+	SPH_C32(0x3C6EF372), SPH_C32(0xA54FF53A),
+	SPH_C32(0x510E527F), SPH_C32(0x9B05688C),
+	SPH_C32(0x1F83D9AB), SPH_C32(0x5BE0CD19)
+};
+
+__device__ __constant__
+static const uint32_t __align__(32) c_u256[16] = {
+	SPH_C32(0x243F6A88), SPH_C32(0x85A308D3),
+	SPH_C32(0x13198A2E), SPH_C32(0x03707344),
+	SPH_C32(0xA4093822), SPH_C32(0x299F31D0),
+	SPH_C32(0x082EFA98), SPH_C32(0xEC4E6C89),
+	SPH_C32(0x452821E6), SPH_C32(0x38D01377),
+	SPH_C32(0xBE5466CF), SPH_C32(0x34E90C6C),
+	SPH_C32(0xC0AC29B7), SPH_C32(0xC97C50DD),
+	SPH_C32(0x3F84D5B5), SPH_C32(0xB5470917)
+};
+
+#if 0
+#define GS(m0, m1, c0, c1, a, b, c, d)   do { \
+		a = SPH_T32(a + b + (m0 ^ c1)); \
+		d = SPH_ROTR32(d ^ a, 16); \
+		c = SPH_T32(c + d); \
+		b = SPH_ROTR32(b ^ c, 12); \
+		a = SPH_T32(a + b + (m1 ^ c0)); \
+		d = SPH_ROTR32(d ^ a, 8); \
+		c = SPH_T32(c + d); \
+		b = SPH_ROTR32(b ^ c, 7); \
+	} while (0)
+
+#define ROUND_S(r)   do { \
+	GS(Mx(r, 0x0), Mx(r, 0x1), CSx(r, 0x0), CSx(r, 0x1), v[0], v[4], v[0x8], v[0xC]); \
+	GS(Mx(r, 0x2), Mx(r, 0x3), CSx(r, 0x2), CSx(r, 0x3), v[1], v[5], v[0x9], v[0xD]); \
+	GS(Mx(r, 0x4), Mx(r, 0x5), CSx(r, 0x4), CSx(r, 0x5), v[2], v[6], v[0xA], v[0xE]); \
+	GS(Mx(r, 0x6), Mx(r, 0x7), CSx(r, 0x6), CSx(r, 0x7), v[3], v[7], v[0xB], v[0xF]); \
+	GS(Mx(r, 0x8), Mx(r, 0x9), CSx(r, 0x8), CSx(r, 0x9), v[0], v[5], v[0xA], v[0xF]); \
+	GS(Mx(r, 0xA), Mx(r, 0xB), CSx(r, 0xA), CSx(r, 0xB), v[1], v[6], v[0xB], v[0xC]); \
+	GS(Mx(r, 0xC), Mx(r, 0xD), CSx(r, 0xC), CSx(r, 0xD), v[2], v[7], v[0x8], v[0xD]); \
+	GS(Mx(r, 0xE), Mx(r, 0xF), CSx(r, 0xE), CSx(r, 0xF), v[3], v[4], v[0x9], v[0xE]); \
+} while (0)
+#endif
+
+#define GS(a,b,c,d,x) { \
+	const uint32_t idx1 = c_sigma[i][x]; \
+	const uint32_t idx2 = c_sigma[i][x+1]; \
+	v[a] += (m[idx1] ^ u256[idx2]) + v[b]; \
+	v[d] = SPH_ROTL32(v[d] ^ v[a], 16); \
+	v[c] += v[d]; \
+	v[b] = SPH_ROTR32(v[b] ^ v[c], 12); \
+\
+	v[a] += (m[idx2] ^ u256[idx1]) + v[b]; \
+	v[d] = SPH_ROTR32(v[d] ^ v[a], 8); \
+	v[c] += v[d]; \
+	v[b] = SPH_ROTR32(v[b] ^ v[c], 7); \
+}
+
+#define BLAKE256_ROUNDS 14
+
+__device__ static
+void blake256_compress(uint32_t *h, const uint32_t *block, const uint32_t T0)
+{
+	uint32_t /* __align__(8) */ v[16];
+	uint32_t /* __align__(8) */ m[16];
+
+	const uint32_t* u256 = c_u256;
+
+	//#pragma unroll
+	for (int i = 0; i < 16; ++i) {
+		m[i] = block[i];
+	}
+
+	//#pragma unroll 8
+	for(int i = 0; i < 8; i++)
+		v[i] = h[i];
+
+	v[ 8] = u256[0];
+	v[ 9] = u256[1];
+	v[10] = u256[2];
+	v[11] = u256[3];
+
+	v[12] = u256[4] ^ T0;
+	v[13] = u256[5] ^ T0;
+	v[14] = u256[6];
+	v[15] = u256[7];
+
+	//#pragma unroll
+	for (int i = 0; i < BLAKE256_ROUNDS; i++) {
+		/* column step */
+		GS(0, 4, 0x8, 0xC, 0);
+		GS(1, 5, 0x9, 0xD, 2);
+		GS(2, 6, 0xA, 0xE, 4);
+		GS(3, 7, 0xB, 0xF, 6);
+		/* diagonal step */
+		GS(0, 5, 0xA, 0xF, 0x8);
+		GS(1, 6, 0xB, 0xC, 0xA);
+		GS(2, 7, 0x8, 0xD, 0xC);
+		GS(3, 4, 0x9, 0xE, 0xE);
+	}
+
+	//#pragma unroll 16
+	for(int i = 0; i < 16; i++)
+		h[i % 8] ^= v[i];
+}
+
+__global__
+void blake256_gpu_hash_80(uint32_t threads, uint32_t startNounce, uint32_t *resNounce)
+{
+	uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
+	if (thread < threads)
+	{
+		const uint32_t nounce = startNounce + thread;
+		uint32_t /* __align__(8) */ msg[16];
+		uint32_t h[8];
+
+		#pragma unroll
+		for(int i=0; i<8; i++)
+			h[i] = c_IV256[i];
+
+		blake256_compress(h, c_PaddedMessage80, 0x200); /* 512 = 0x200 */
+
+		// ------ Close: Bytes 64 to 80 ------ 
+
+		msg[0] = c_PaddedMessage80[16];
+		msg[1] = c_PaddedMessage80[17];
+		msg[2] = c_PaddedMessage80[18];
+		msg[3] = nounce; /* our tested value */
+		msg[4] = 0x80000000UL; //cuda_swab32(0x80U);
+
+		msg[5] = 0;  // uchar[17 to 55]
+		msg[6] = 0;
+		msg[7] = 0;
+		msg[8] = 0;
+		msg[9] = 0;
+		msg[10] = 0;
+		msg[11] = 0;
+		msg[12] = 0;
+
+		msg[13] = 1;
+		msg[14] = 0;
+		msg[15] = 0x280;
+
+		blake256_compress(h, msg, 0x280);
+
+		for (int i = 7; i >= 0; i--) {
+			uint32_t hash = cuda_swab32(h[i]);
+			if (hash > c_Target[i]) {
+				return;
+			}
+			if (hash < c_Target[i]) {
+				break;
+			}
+		}
+
+		/* keep the smallest nounce, hmm... */
+		if(resNounce[0] > nounce)
+			resNounce[0] = nounce;
+	}
+}
+
+__host__
+uint32_t blake256_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce)
+{
+	const int threadsperblock = TPB;
+	uint32_t result = MAXU;
+
+	dim3 grid((threads + threadsperblock-1)/threadsperblock);
+	dim3 block(threadsperblock);
+	size_t shared_size = 0;
+
+	/* Check error on Ctrl+C or kill to prevent segfaults on exit */
+	if (cudaMemset(d_resNounce[thr_id], 0xff, sizeof(uint32_t)) != cudaSuccess)
+		return result;
+
+	blake256_gpu_hash_80<<<grid, block, shared_size>>>(threads, startNounce, d_resNounce[thr_id]);
+	cudaDeviceSynchronize();
+	if (cudaSuccess == cudaMemcpy(h_resNounce[thr_id], d_resNounce[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost)) {
+		cudaThreadSynchronize();
+		result = *h_resNounce[thr_id];
+	}
+	return result;
+}
+
+__host__
+void blake256_cpu_setBlock_80(uint32_t *pdata, const uint32_t *ptarget)
+{
+	uint32_t PaddedMessage[32];
+	memcpy(PaddedMessage, pdata, 80);
+	memset(&PaddedMessage[20], 0, 48);
+	CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_PaddedMessage80, PaddedMessage, sizeof(PaddedMessage), 0, cudaMemcpyHostToDevice));
+	CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_sigma, host_sigma, sizeof(host_sigma), 0, cudaMemcpyHostToDevice));
+	CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_Target, ptarget, 32, 0, cudaMemcpyHostToDevice));
+}
+
+extern "C" int scanhash_blake32(int thr_id, uint32_t *pdata, const uint32_t *ptarget,
+	uint32_t max_nonce, unsigned long *hashes_done)
+{
+	const uint32_t first_nonce = pdata[19];
+	static bool init[8] = { 0, 0, 0, 0, 0, 0, 0, 0 };
+	uint32_t throughput = min(TPB * 2048, max_nonce - first_nonce);
+	int rc = 0;
+
+	if (opt_benchmark)
+		((uint32_t*)ptarget)[7] = 0x00000f;
+
+	if (!init[thr_id]) {
+		if (opt_n_threads > 1) {
+			CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id]));
+		}
+		CUDA_SAFE_CALL(cudaMallocHost(&h_resNounce[thr_id], sizeof(uint32_t)));
+		CUDA_SAFE_CALL(cudaMalloc(&d_resNounce[thr_id], sizeof(uint32_t)));
+		init[thr_id] = true;
+	}
+
+	if (throughput < (TPB * 2048))
+		applog(LOG_WARNING, "throughput=%u, start=%x, max=%x", throughput, first_nonce, max_nonce);
+
+	blake256_cpu_setBlock_80(pdata, ptarget);
+
+	do {
+		// GPU HASH
+		uint32_t foundNonce = blake256_cpu_hash_80(thr_id, throughput, pdata[19]);
+		if (foundNonce != 0xffffffff)
+		{
+			uint32_t endiandata[20];
+			uint32_t vhashcpu[8];
+			uint32_t Htarg = ptarget[7];
+
+			for (int k=0; k < 20; k++)
+				be32enc(&endiandata[k], pdata[k]);
+
+			if (opt_debug && !opt_quiet) {
+				applog(LOG_DEBUG, "throughput=%u, start=%x, max=%x, pdata=%08x...%08x",
+					throughput, first_nonce, max_nonce, endiandata[0], endiandata[7]);
+				applog_hash((unsigned char *)pdata);
+			}
+
+			be32enc(&endiandata[19], foundNonce);
+
+			blake32hash(vhashcpu, endiandata);
+
+			if (vhashcpu[7] <= Htarg && fulltest(vhashcpu, ptarget))
+			{
+				pdata[19] = foundNonce;
+				rc = 1;
+				goto exit_scan;
+			}
+			else if (vhashcpu[7] > Htarg) {
+				applog(LOG_WARNING, "GPU #%d: result for nounce %08x is not in range: %x > %x", thr_id, foundNonce, vhashcpu[7], Htarg);
+			}
+			else if (vhashcpu[6] > ptarget[6]) {
+				applog(LOG_WARNING, "GPU #%d: hash[6] for nounce %08x is not in range: %x > %x", thr_id, foundNonce, vhashcpu[6], ptarget[6]);
+			}
+			else {
+				applog(LOG_WARNING, "GPU #%d: result for nounce %08x does not validate on CPU!", thr_id, foundNonce);
+			}
+		}
+
+		pdata[19] += throughput;
+
+	} while (pdata[19] < max_nonce && !work_restart[thr_id].restart);
+
+exit_scan:
+	*hashes_done = pdata[19] - first_nonce + 1;
+	// reset the device to allow multiple instances
+	if (opt_n_threads == 1) {
+		CUDA_SAFE_CALL(cudaDeviceReset());
+		init[thr_id] = false;
+	}
+	// wait proper end of all threads
+	cudaDeviceSynchronize();
+	return rc;
+}

ccminer.vcxproj

@@ -243,6 +243,7 @@ copy "$(CudaToolkitBinDir)\cudart64*.dll" "$(OutDir)"</Command>
     </ClCompile>
     <ClCompile Include="fuguecoin.cpp" />
     <ClCompile Include="groestlcoin.cpp" />
+    <ClCompile Include="hashlog.cpp" />
     <ClCompile Include="hefty1.c" />
     <ClCompile Include="myriadgroestl.cpp" />
     <ClCompile Include="scrypt.c">
@@ -397,6 +398,11 @@ copy "$(CudaToolkitBinDir)\cudart64*.dll" "$(OutDir)"</Command>
       <AdditionalOptions Condition="'$(Configuration)'=='Debug'">%(AdditionalOptions)</AdditionalOptions>
       <TargetMachinePlatform Condition="'$(Platform)'=='x64'">64</TargetMachinePlatform>
     </CudaCompile>
+    <CudaCompile Include="blake32.cu">
+      <MaxRegCount>64</MaxRegCount>
+      <AdditionalOptions Condition="'$(Configuration)'=='Release'">--ptxas-options=-O2 %(AdditionalOptions)</AdditionalOptions>
+      <AdditionalOptions Condition="'$(Configuration)'=='Debug'">%(AdditionalOptions)</AdditionalOptions>
+    </CudaCompile>
     <CudaCompile Include="quark\animecoin.cu">
       <AdditionalOptions Condition="'$(Configuration)'=='Release'">--ptxas-options=-O2 %(AdditionalOptions)</AdditionalOptions>
       <AdditionalOptions Condition="'$(Configuration)'=='Debug'">%(AdditionalOptions)</AdditionalOptions>

ccminer.vcxproj.filters

@@ -180,6 +180,9 @@
     <ClCompile Include="compat\winansi.c">
       <Filter>Source Files</Filter>
     </ClCompile>
+    <ClCompile Include="hashlog.cpp">
+      <Filter>Source Files</Filter>
+    </ClCompile>
   </ItemGroup>
   <ItemGroup>
     <ClInclude Include="compat.h">
@@ -436,5 +439,8 @@
     <CudaCompile Include="x17\x17.cu">
       <Filter>Source Files\CUDA\x17</Filter>
     </CudaCompile>
+    <CudaCompile Include="blake32.cu">
+      <Filter>Source Files\CUDA</Filter>
+    </CudaCompile>
   </ItemGroup>
 </Project>

configure.ac

@@ -1,4 +1,4 @@
-AC_INIT([ccminer], [2014.08.12])
+AC_INIT([ccminer], [2014.09.01])
 
 AC_PREREQ([2.59c])
 AC_CANONICAL_SYSTEM
@@ -144,12 +144,12 @@ AC_ARG_WITH([cuda],
 
 if test -n "$with_cuda"
 then
-   CUDA_CFLAGS="-I$with_cuda/include"
+   CUDA_CFLAGS="-I$with_cuda/include $CUDA_CFLAGS"
    CUDA_LIBS="-lcudart"
    CUDA_LDFLAGS="-L$with_cuda/lib$SUFFIX"
    NVCC="$with_cuda/bin/nvcc"
 else
-   CUDA_CFLAGS="-I/usr/local/cuda/include"
+   CUDA_CFLAGS="-I/usr/local/cuda/include $CUDA_CFLAGS"
    CUDA_LIBS="-lcudart -static-libstdc++"
    CUDA_LDFLAGS="-L/usr/local/cuda/lib$SUFFIX"
    NVCC="nvcc"

configure.sh

@@ -1 +1,9 @@
-./configure "CFLAGS=-O3" "CXXFLAGS=-O3" --with-cuda=/usr/local/cuda
+# possible additional CUDA_CFLAGS
+#-gencode=arch=compute_50,code=\"sm_50,compute_50\"
+#-gencode=arch=compute_35,code=\"sm_35,compute_35\"
+#-gencode=arch=compute_30,code=\"sm_30,compute_30\"
+
+#--ptxas-options=\"-v -dlcm=cg\""
+
+CUDA_CFLAGS="-O2" ./configure "CFLAGS=-O2" "CXXFLAGS=-O2" --with-cuda=/usr/local/cuda
+

cpu-miner.c

@@ -47,7 +47,7 @@ BOOL WINAPI ConsoleHandler(DWORD);
 #pragma comment(lib, "winmm.lib")
 #endif
 
-#define PROGRAM_NAME		"minerd"
+#define PROGRAM_NAME		"ccminer"
 #define LP_SCANTIME		60
 #define HEAVYCOIN_BLKHDR_SZ		84
 #define MNR_BLKHDR_SZ 80
@@ -126,15 +126,16 @@ struct workio_cmd {
 };
 
 typedef enum {
-	ALGO_HEAVY,		/* Heavycoin hash */
+	ALGO_ANIME,
-	ALGO_MJOLLNIR,		/* Mjollnir hash */
+	ALGO_BLAKE,
+	ALGO_FRESH,
 	ALGO_FUGUE256,		/* Fugue256 */
 	ALGO_GROESTL,
-	ALGO_MYR_GR,
+	ALGO_HEAVY,		/* Heavycoin hash */
 	ALGO_JACKPOT,
+	ALGO_MJOLLNIR,		/* Mjollnir hash */
+	ALGO_MYR_GR,
 	ALGO_QUARK,
-	ALGO_ANIME,
-	ALGO_FRESH,
 	ALGO_NIST5,
 	ALGO_WHC,
 	ALGO_X11,
@@ -146,16 +147,17 @@ typedef enum {
 } sha256_algos;
 
 static const char *algo_names[] = {
-	"heavy",
+	"anime",
-	"mjollnir",
+	"blake",
+	"fresh",
 	"fugue256",
 	"groestl",
-	"myr-gr",
+	"heavy",
 	"jackpot",
-	"quark",
+	"mjollnir",
-	"anime",
+	"myr-gr",
-	"fresh",
 	"nist5",
+	"quark",
 	"whirl",
 	"x11",
 	"x13",
@@ -166,6 +168,7 @@ static const char *algo_names[] = {
 };
 
 bool opt_debug = false;
+bool opt_debug_rpc = false;
 bool opt_protocol = false;
 bool opt_benchmark = false;
 bool want_longpoll = true;
@@ -176,7 +179,7 @@ static bool submit_old = false;
 bool use_syslog = false;
 bool use_colors = false;
 static bool opt_background = false;
-static bool opt_quiet = false;
+bool opt_quiet = false;
 static int opt_retries = -1;
 static int opt_fail_pause = 30;
 int opt_timeout = 270;
@@ -184,7 +187,7 @@ static int opt_scantime = 5;
 static json_t *opt_config;
 static const bool opt_time = true;
 static sha256_algos opt_algo = ALGO_HEAVY;
-static int opt_n_threads = 0;
+int opt_n_threads = 0;
 static double opt_difficulty = 1; // CH
 bool opt_trust_pool = false;
 uint16_t opt_vote = 9999;
@@ -207,7 +210,6 @@ static struct stratum_ctx stratum;
 
 pthread_mutex_t applog_lock;
 static pthread_mutex_t stats_lock;
-
 static unsigned long accepted_count = 0L;
 static unsigned long rejected_count = 0L;
 static double *thr_hashrates;
@@ -227,16 +229,17 @@ static char const usage[] = "\
 Usage: " PROGRAM_NAME " [OPTIONS]\n\
 Options:\n\
   -a, --algo=ALGO       specify the algorithm to use\n\
-                        fugue256  Fuguecoin hash\n\
-                        heavy     Heavycoin hash\n\
-                        mjollnir  Mjollnircoin hash\n\
-                        groestl   Groestlcoin hash\n\
-                        myr-gr    Myriad-Groestl hash\n\
-                        jackpot   Jackpot hash\n\
-                        quark     Quark hash\n\
                         anime     Animecoin hash\n\
+                        blake     Blake 256 (like NEOS blake)\n\
                         fresh     Freshcoin hash (shavite 80)\n\
+                        fugue256  Fuguecoin hash\n\
+                        groestl   Groestlcoin hash\n\
+                        heavy     Heavycoin hash\n\
+                        jackpot   Jackpot hash\n\
+                        mjollnir  Mjollnircoin hash\n\
+                        myr-gr    Myriad-Groestl hash\n\
                         nist5     NIST5 (TalkCoin) hash\n\
+                        quark     Quark hash\n\
                         whirl     Whirlcoin (old whirlpool)\n\
                         x11       X11 (DarkCoin) hash\n\
                         x13       X13 (MaruCoin) hash\n\
@@ -340,12 +343,23 @@ struct work {
 	char job_id[128];
 	size_t xnonce2_len;
 	unsigned char xnonce2[32];
+
+	uint32_t scanned_from;
+	uint32_t scanned_to;
 };
 
 static struct work g_work;
 static time_t g_work_time;
 static pthread_mutex_t g_work_lock;
 
+
+void proper_exit(int reason)
+{
+	cuda_devicereset();
+	hashlog_purge_all();
+	exit(reason);
+}
+
 static bool jobj_binary(const json_t *obj, const char *key,
 			void *buf, size_t buflen)
 {
@@ -397,11 +411,11 @@ err_out:
 	return false;
 }
 
-static void share_result(int result, const char *reason)
+static int share_result(int result, const char *reason)
 {
 	char s[345];
 	double hashrate;
-	int i;
+	int i, ret = 0;
 
 	hashrate = 0.;
 	pthread_mutex_lock(&stats_lock);
@@ -417,11 +431,18 @@ static void share_result(int result, const char *reason)
 			100. * accepted_count / (accepted_count + rejected_count),
 			s,
 			use_colors ?
-				(result ? CL_GRN "(yay!!!)" : CL_RED "(booooo)")
+				(result ? CL_GRN "yay!!!" : CL_RED "booooo")
 			:	(result ? "(yay!!!)" : "(booooo)"));
 
-	if (opt_debug && reason)
+	if (reason && !opt_quiet) {
-		applog(LOG_DEBUG, "DEBUG: reject reason: %s", reason);
+		applog(LOG_WARNING, "reject reason: %s", reason);
+		if (strncmp(reason, "low difficulty share", 20) == 0) {
+			opt_difficulty = (opt_difficulty * 2.0) / 3.0;
+			applog(LOG_WARNING, "factor reduced to : %0.2f", opt_difficulty);
+			return 0;
+		}
+	}
+	return 1;
 }
 
 static bool submit_upstream_work(CURL *curl, struct work *work)
@@ -440,6 +461,7 @@ static bool submit_upstream_work(CURL *curl, struct work *work)
 	}
 
 	if (have_stratum) {
+		uint32_t sent;
 		uint32_t ntime, nonce;
 		uint16_t nvote;
 		char *ntimestr, *noncestr, *xnonce2str, *nvotestr;
@@ -452,14 +474,26 @@ static bool submit_upstream_work(CURL *curl, struct work *work)
 		noncestr = bin2hex((const unsigned char *)(&nonce), 4);
 		xnonce2str = bin2hex(work->xnonce2, work->xnonce2_len);
 		nvotestr = bin2hex((const unsigned char *)(&nvote), 2);
+
+		sent = hashlog_already_submittted(work->job_id, nonce);
+		if (sent > 0) {
+			sent = (uint32_t) time(NULL) - sent;
+			if (!opt_quiet) {
+				applog(LOG_WARNING, "skip submit, nonce %s was already sent %u seconds ago", noncestr, sent);
+				hashlog_dump_job(work->job_id);
+			}
+			rc = true;
+			goto out;
+		}
+
 		if (opt_algo == ALGO_HEAVY) {
 			sprintf(s,
 				"{\"method\": \"mining.submit\", \"params\": [\"%s\", \"%s\", \"%s\", \"%s\", \"%s\", \"%s\"], \"id\":4}",
-				rpc_user, work->job_id, xnonce2str, ntimestr, noncestr, nvotestr);
+				rpc_user, work->job_id + 8, xnonce2str, ntimestr, noncestr, nvotestr);
 		} else {
 			sprintf(s,
 				"{\"method\": \"mining.submit\", \"params\": [\"%s\", \"%s\", \"%s\", \"%s\", \"%s\"], \"id\":4}",
-				rpc_user, work->job_id, xnonce2str, ntimestr, noncestr);
+				rpc_user, work->job_id + 8, xnonce2str, ntimestr, noncestr);
 		}
 		free(ntimestr);
 		free(noncestr);
@@ -470,6 +504,10 @@ static bool submit_upstream_work(CURL *curl, struct work *work)
 			applog(LOG_ERR, "submit_upstream_work stratum_send_line failed");
 			goto out;
 		}
+
+		hashlog_remember_submit(work->job_id, nonce);
+		hashlog_remember_scan_range(work->job_id, work->scanned_from, work->scanned_to);
+
 	} else {
 
 		/* build hex string */
@@ -498,11 +536,16 @@ static bool submit_upstream_work(CURL *curl, struct work *work)
 
 		res = json_object_get(val, "result");
 		reason = json_object_get(val, "reject-reason");
-		share_result(json_is_true(res), reason ? json_string_value(reason) : NULL);
+		if (!share_result(json_is_true(res), reason ? json_string_value(reason) : NULL))
+			hashlog_purge_job(work->job_id);
 
 		json_decref(val);
 	}
 
+	if (opt_debug_rpc) {
+		applog(LOG_DEBUG, "submit: %s", s);
+	}
+
 	rc = true;
 
 out:
@@ -734,7 +777,9 @@ static void stratum_gen_work(struct stratum_ctx *sctx, struct work *work)
 
 	pthread_mutex_lock(&sctx->work_lock);
 
-	strcpy(work->job_id, sctx->job.job_id);
+	// store the job ntime as high part of jobid
+	snprintf(work->job_id, sizeof(work->job_id), "%07x %s",
+		be32dec(sctx->job.ntime) & 0xfffffff, sctx->job.job_id);
 	work->xnonce2_len = sctx->xnonce2_size;
 	memcpy(work->xnonce2, sctx->job.xnonce2, sctx->xnonce2_size);
 
@@ -759,7 +804,7 @@ static void stratum_gen_work(struct stratum_ctx *sctx, struct work *work)
 	for (i = 0; i < (int)sctx->xnonce2_size && !++sctx->job.xnonce2[i]; i++);
 
 	/* Assemble block header */
-	memset(work->data, 0, 128);
+	memset(work->data, 0, sizeof(work->data));
 	work->data[0] = le32dec(sctx->job.version);
 	for (i = 0; i < 8; i++)
 		work->data[1 + i] = le32dec((uint32_t *)sctx->job.prevhash + i);
@@ -792,9 +837,11 @@ static void stratum_gen_work(struct stratum_ctx *sctx, struct work *work)
 	pthread_mutex_unlock(&sctx->work_lock);
 
 	if (opt_debug) {
+		char *tm = atime2str(swab32(work->data[17]) - sctx->srvtime_diff);
 		char *xnonce2str = bin2hex(work->xnonce2, sctx->xnonce2_size);
-		applog(LOG_DEBUG, "DEBUG: job_id='%s' extranonce2=%s ntime=%08x",
+		applog(LOG_DEBUG, "DEBUG: job_id=%s xnonce2=%s time=%s",
-		       work->job_id, xnonce2str, swab32(work->data[17]));
+		       work->job_id, xnonce2str, tm);
+		free(tm);
 		free(xnonce2str);
 	}
 
@@ -812,10 +859,10 @@ static void *miner_thread(void *userdata)
 	int thr_id = mythr->id;
 	struct work work;
 	uint32_t max_nonce;
-	uint32_t end_nonce = 0xffffffffU / opt_n_threads * (thr_id + 1) - 0x20;
+	uint32_t end_nonce = 0xffffffffU / opt_n_threads * (thr_id + 1) - (thr_id + 1);
 	unsigned char *scratchbuf = NULL;
+	bool work_done = false;
 	char s[16];
-	int i;
 
 	memset(&work, 0, sizeof(work)); // prevent work from being used uninitialized
 
@@ -838,56 +885,138 @@ static void *miner_thread(void *userdata)
 
 	while (1) {
 		unsigned long hashes_done;
+		uint32_t start_nonce;
 		struct timeval tv_start, tv_end, diff;
 		int64_t max64;
+		uint64_t umax64;
 		int rc;
 
+		// &work.data[19]
+		int wcmplen = 76;
+		uint32_t *nonceptr = (uint32_t*) (((char*)work.data) + wcmplen);
+
 		if (have_stratum) {
-			while (time(NULL) >= g_work_time + 120)
+			while (time(NULL) >= (g_work_time + opt_scantime) && !work_done)
-				sleep(1);
+				usleep(500*1000);
+			work_done = false;
 			pthread_mutex_lock(&g_work_lock);
-			if (work.data[19] >= end_nonce)
+			nonceptr = (uint32_t*) (((char*)work.data) + wcmplen);
+			if ((*nonceptr) >= end_nonce)
 				stratum_gen_work(&stratum, &g_work);
 		} else {
+			int min_scantime = have_longpoll ? LP_SCANTIME : opt_scantime;
 			/* obtain new work from internal workio thread */
 			pthread_mutex_lock(&g_work_lock);
-			if (!have_stratum && (!have_longpoll ||
+			if (time(NULL) - g_work_time >= min_scantime ||
-					time(NULL) >= g_work_time + LP_SCANTIME*3/4 ||
+			     (*nonceptr) >= end_nonce) {
-					work.data[19] >= end_nonce)) {
 				if (unlikely(!get_work(mythr, &g_work))) {
 					applog(LOG_ERR, "work retrieval failed, exiting "
 						"mining thread %d", mythr->id);
 					pthread_mutex_unlock(&g_work_lock);
 					goto out;
 				}
-				g_work_time = have_stratum ? 0 : time(NULL);
+				g_work_time = time(NULL);
-			}
+			}
-			if (have_stratum) {
+		}
-				pthread_mutex_unlock(&g_work_lock);
+		if (memcmp(work.data, g_work.data, wcmplen)) {
-				continue;
+			if (opt_debug) {
+				applog(LOG_DEBUG, "job %s work updated", g_work.job_id);
+				for (int n=0; n<wcmplen; n+=8) {
+					if (memcmp(work.data + n, g_work.data + n, 8)) {
+						applog(LOG_DEBUG, "diff detected at offset %d", n);
+						applog_hash((uint8_t*) work.data + n);
+						applog_hash((uint8_t*) g_work.data + n);
+					}
 				}
 			}
-		if (memcmp(work.data, g_work.data, 76)) {
 			memcpy(&work, &g_work, sizeof(struct work));
-			work.data[19] = 0xffffffffU / opt_n_threads * thr_id;
+			(*nonceptr) = (0xffffffffUL / opt_n_threads) * thr_id; // 0 if single thr
+		} else if (memcmp(work.target, g_work.target, sizeof(work.target))) {
+			if (opt_debug) {
+				applog(LOG_DEBUG, "job %s target change", g_work.job_id);
+				applog_hash((uint8_t*) work.target);
+				applog_hash((uint8_t*) g_work.target);
+			}
+			memcpy(work.target, g_work.target, sizeof(work.target));
+			(*nonceptr) = (0xffffffffUL / opt_n_threads) * thr_id; // 0 if single thr
 		} else
-			work.data[19]++;
+			(*nonceptr)++; //??
 		pthread_mutex_unlock(&g_work_lock);
 		work_restart[thr_id].restart = 0;
 
+		if (opt_debug)
+			applog(LOG_WARNING, "job %s %08x", g_work.job_id, (*nonceptr));
+
 		/* adjust max_nonce to meet target scan time */
 		if (have_stratum)
 			max64 = LP_SCANTIME;
 		else
 			max64 = g_work_time + (have_longpoll ? LP_SCANTIME : opt_scantime)
 			      - time(NULL);
+
 		max64 *= (int64_t)thr_hashrates[thr_id];
-		if (max64 <= 0)
+
-			max64 = (opt_algo == ALGO_JACKPOT) ? 0x1fffLL : 0xfffffLL;
+		if (max64 <= 0) {
-		if ((int64_t)work.data[19] + max64 > end_nonce)
+			switch (opt_algo) {
+			case ALGO_JACKPOT:
+				max64 = 0x1fffLL;
+				break;
+			case ALGO_BLAKE:
+				/* based on the 750Ti hashrate (100kH) */
+				max64 = 0x3ffffffLL;
+				break;
+			default:
+				max64 = 0xfffffLL;
+				break;
+			}
+		}
+
+		start_nonce = *nonceptr;
+
+		/* do not recompute something already scanned */
+		if (opt_algo == ALGO_BLAKE && opt_n_threads == 1) {
+			union {
+				uint64_t data;
+				uint32_t scanned[2];
+			} range;
+
+			range.data = hashlog_get_scan_range(work.job_id);
+			if (range.data) {
+				bool stall = false;
+				if (range.scanned[0] == 1 && range.scanned[1] == 0xFFFFFFFFUL) {
+					applog(LOG_WARNING, "detected a rescan of fully scanned job!");
+				} else if (range.scanned[0] > 0 && range.scanned[1] > 0 && range.scanned[1] < 0xFFFFFFF0UL) {
+					/* continue scan the end */
+					start_nonce = range.scanned[1] + 1;
+					//applog(LOG_DEBUG, "scan the next part %x + 1 (%x-%x)", range.scanned[1], range.scanned[0], range.scanned[1]);
+				}
+
+				stall = (start_nonce == work.scanned_from && end_nonce == work.scanned_to);
+				stall |= (start_nonce == work.scanned_from && start_nonce == range.scanned[1] + 1);
+				stall |= (start_nonce > range.scanned[0] && start_nonce < range.scanned[1]);
+
+				if (stall) {
+					if (opt_algo)
+						applog(LOG_DEBUG, "job done, wait for a new one...");
+					work_restart[thr_id].restart = 1;
+					hashlog_purge_old();
+					// wait a bit for a new job...
+					usleep(1500*1000);
+					(*nonceptr) = end_nonce + 1;
+					work_done = true;
+					continue;
+				}
+			}
+		}
+
+		umax64 = (uint64_t) max64;
+		if ((umax64 + start_nonce) >= end_nonce)
 			max_nonce = end_nonce;
 		else
-			max_nonce = (uint32_t)(work.data[19] + max64);
+			max_nonce = (uint32_t) umax64 + start_nonce;
+
+		work.scanned_from = start_nonce;
+		(*nonceptr) = start_nonce;
 
 		hashes_done = 0;
 		gettimeofday(&tv_start, NULL);
@@ -936,6 +1065,11 @@ static void *miner_thread(void *userdata)
 			                      max_nonce, &hashes_done);
 			break;
 
+		case ALGO_BLAKE:
+			rc = scanhash_blake32(thr_id, work.data, work.target,
+			                      max_nonce, &hashes_done);
+			break;
+
 		case ALGO_FRESH:
 			rc = scanhash_fresh(thr_id, work.data, work.target,
 			                      max_nonce, &hashes_done);
@@ -983,6 +1117,10 @@ static void *miner_thread(void *userdata)
 
 		/* record scanhash elapsed time */
 		gettimeofday(&tv_end, NULL);
+
+		if (rc && opt_debug)
+			applog(LOG_NOTICE, CL_CYN "found => %08x" CL_GRN " %08x", *nonceptr, swab32(*nonceptr));
+
 		timeval_subtract(&diff, &tv_end, &tv_start);
 		if (diff.tv_usec || diff.tv_sec) {
 			pthread_mutex_lock(&stats_lock);
@@ -993,19 +1131,29 @@ static void *miner_thread(void *userdata)
 		if (!opt_quiet) {
 			sprintf(s, thr_hashrates[thr_id] >= 1e6 ? "%.0f" : "%.2f",
 				1e-3 * thr_hashrates[thr_id]);
-			applog(LOG_INFO, "GPU #%d: %s, %s khash/s",
+			applog(LOG_INFO, "GPU #%d: %s, %s kH/s",
 				device_map[thr_id], device_name[thr_id], s);
 		}
 		if (opt_benchmark && thr_id == opt_n_threads - 1) {
 			double hashrate = 0.;
+			int i;
 			for (i = 0; i < opt_n_threads && thr_hashrates[i]; i++)
 				hashrate += thr_hashrates[i];
 			if (i == opt_n_threads) {
 				sprintf(s, hashrate >= 1e6 ? "%.0f" : "%.2f", hashrate / 1000.);
-				applog(LOG_NOTICE, "Total: %s khash/s", s);
+				applog(LOG_NOTICE, "Total: %s kH/s", s);
 			}
 		}
 
+		if (rc) {
+			work.scanned_to = *nonceptr;
+		} else {
+			work.scanned_to = max_nonce;
+		}
+
+		// could be used to store speeds too..
+		hashlog_remember_scan_range(work.job_id, work.scanned_from, work.scanned_to);
+
 		/* if nonce found, submit work */
 		if (rc && !opt_benchmark && !submit_work(mythr, &work))
 			break;
@@ -1181,15 +1329,20 @@ static void *stratum_thread(void *userdata)
 		}
 
 		if (stratum.job.job_id &&
-		    (strcmp(stratum.job.job_id, g_work.job_id) || !g_work_time)) {
+		    (!g_work_time || strncmp(stratum.job.job_id, g_work.job_id + 8, 120))) {
 			pthread_mutex_lock(&g_work_lock);
 			stratum_gen_work(&stratum, &g_work);
 			time(&g_work_time);
 			pthread_mutex_unlock(&g_work_lock);
 			if (stratum.job.clean) {
 				if (!opt_quiet)
-					applog(LOG_BLUE, "%s send a new %s block", short_url, algo_names[opt_algo]);
+					applog(LOG_BLUE, "%s requested %s job %d restart, block %d", short_url, algo_names[opt_algo],
+						strtoul(stratum.job.job_id, NULL, 16), stratum.bloc_height);
 				restart_threads();
+				hashlog_purge_old();
+			} else if (!opt_quiet) {
+					applog(LOG_BLUE, "%s send %s job %d, block %d", short_url, algo_names[opt_algo],
+						strtoul(stratum.job.job_id, NULL, 16), stratum.bloc_height);
 			}
 		}
 		
@@ -1212,7 +1365,7 @@ out:
 	return NULL;
 }
 
-#define PROGRAM_VERSION "1.3"
+#define PROGRAM_VERSION "1.4"
 static void show_version_and_exit(void)
 {
 	printf("%s v%s\n"
@@ -1225,7 +1378,7 @@ static void show_version_and_exit(void)
 		PTW32_VERSION_STRING,
 #endif
 		curl_version());
-	exit(0);
+	proper_exit(0);
 }
 
 static void show_usage_and_exit(int status)
@@ -1234,7 +1387,7 @@ static void show_usage_and_exit(int status)
 		fprintf(stderr, "Try `" PROGRAM_NAME " --help' for more information.\n");
 	else
 		printf(usage);
-	exit(status);
+	proper_exit(status);
 }
 
 static void parse_arg (int key, char *arg)
@@ -1269,7 +1422,7 @@ static void parse_arg (int key, char *arg)
 #endif
 		if (!json_is_object(opt_config)) {
 			applog(LOG_ERR, "JSON decode of %s failed", arg);
-			exit(1);
+			proper_exit(1);
 		}
 		break;
 	}
@@ -1281,6 +1434,7 @@ static void parse_arg (int key, char *arg)
 		break;
 	case 'D':
 		opt_debug = true;
+		opt_debug_rpc = true;
 		break;
 	case 'p':
 		free(rpc_pass);
@@ -1412,7 +1566,7 @@ static void parse_arg (int key, char *arg)
 		break;
 	case 1006:
 		print_hash_tests();
-		exit(0);
+		proper_exit(0);
 		break;
 	case 1003:
 		want_longpoll = false;
@@ -1434,7 +1588,7 @@ static void parse_arg (int key, char *arg)
 						device_map[opt_n_threads++] = atoi(pch);
 					else {
 						applog(LOG_ERR, "Non-existant CUDA device #%d specified in -d option", atoi(pch));
-						exit(1);
+						proper_exit(1);
 					}
 				} else {
 					int device = cuda_finddevice(pch);
@@ -1442,7 +1596,7 @@ static void parse_arg (int key, char *arg)
 						device_map[opt_n_threads++] = device;
 					else {
 						applog(LOG_ERR, "Non-existant CUDA device '%s' specified in -d option", pch);
-						exit(1);
+						proper_exit(1);
 					}
 				}
 				pch = strtok (NULL, ",");
@@ -1544,13 +1698,11 @@ static void signal_handler(int sig)
 	case SIGINT:
 		signal(sig, SIG_IGN);
 		applog(LOG_INFO, "SIGINT received, exiting");
-		cuda_devicereset();
+		proper_exit(0);
-		exit(0);
 		break;
 	case SIGTERM:
 		applog(LOG_INFO, "SIGTERM received, exiting");
-		cuda_devicereset();
+		proper_exit(0);
-		exit(0);
 		break;
 	}
 }
@@ -1560,13 +1712,11 @@ BOOL WINAPI ConsoleHandler(DWORD dwType)
 	switch (dwType) {
 	case CTRL_C_EVENT:
 		applog(LOG_INFO, "CTRL_C_EVENT received, exiting");
-		cuda_devicereset();
+		proper_exit(0);
-		exit(0);
 		break;
 	case CTRL_BREAK_EVENT:
 		applog(LOG_INFO, "CTRL_BREAK_EVENT received, exiting");
-		cuda_devicereset();
+		proper_exit(0);
-		exit(0);
 		break;
 	default:
 		return false;
@@ -1592,7 +1742,7 @@ int main(int argc, char *argv[])
 	printf("\t    and HVC extension from http://hvc.1gh.com/" "\n\n");
 	printf("\tCuda additions Copyright 2014 Christian Buchner, Christian H.\n");
 	printf("\t  BTC donation address: 16hJF5mceSojnTD3ZTUDqdRhDyPJzoRakM\n");
-	printf("\tCleaned and optimized by Tanguy Pruvot\n");
+	printf("\tInclude some of djm34 additions, cleaned by Tanguy Pruvot\n");
 	printf("\t  BTC donation address: 1AJdfCpLWPNoAMDfHF1wD5y8VgKSSTHxPo\n\n");
 
 	rpc_user = strdup("");
@@ -1618,6 +1768,9 @@ int main(int argc, char *argv[])
 		sprintf(rpc_userpass, "%s:%s", rpc_user, rpc_pass);
 	}
 
+	/* init stratum data.. */
+	memset(&stratum.url, 0, sizeof(stratum));
+
 	pthread_mutex_init(&stats_lock, NULL);
 	pthread_mutex_init(&g_work_lock, NULL);
 	pthread_mutex_init(&stratum.sock_lock, NULL);
@@ -1757,8 +1910,5 @@ int main(int argc, char *argv[])
 
 	applog(LOG_INFO, "workio thread dead, exiting.");
 
-	// nvprof requires this
-	cuda_devicereset();
-
 	return 0;
 }

cpuminer-config.h

@@ -156,7 +156,7 @@
 #define PACKAGE_NAME "ccminer"
 
 /* Define to the full name and version of this package. */
-#define PACKAGE_STRING "ccminer 2014.08.12"
+#define PACKAGE_STRING "ccminer 2014.09.01"
 
 /* Define to the one symbol short name of this package. */
 #define PACKAGE_TARNAME "ccminer"
@@ -165,7 +165,7 @@
 #define PACKAGE_URL ""
 
 /* Define to the version of this package. */
-#define PACKAGE_VERSION "2014.08.12"
+#define PACKAGE_VERSION "2014.09.01"
 
 /* If using the C implementation of alloca, define if you know the
    direction of stack growth for your system; otherwise it will be
@@ -188,7 +188,7 @@
 #define USE_XOP 1
 
 /* Version number of package */
-#define VERSION "2014.08.12"
+#define VERSION "2014.09.01"
 
 /* Define curl_free() as free() if our version of curl lacks curl_free. */
 /* #undef curl_free */

hashlog.cpp

@@ -0,0 +1,226 @@
+#include <stdlib.h>
+#include <memory.h>
+#include <map>
+
+#include "miner.h"
+
+#define HI_DWORD(u64) ((uint32_t) (u64 >> 32))
+#define LO_DWORD(u64) ((uint32_t) u64)
+#define MK_HI64(u32) (0x100000000ULL * u32)
+
+struct hashlog_data {
+	uint32_t tm_sent;
+	uint32_t scanned_from;
+	uint32_t scanned_to;
+	uint32_t last_from;
+	uint32_t tm_add;
+	uint32_t tm_upd;
+};
+
+static std::map<uint64_t, hashlog_data> tlastshares;
+
+#define LOG_PURGE_TIMEOUT 5*60
+
+/**
+ * str hex to uint32
+ */
+static uint64_t hextouint(char* jobid)
+{
+	char *ptr;
+	return strtoull(jobid, &ptr, 16);
+}
+
+/**
+ * @return time of a job/nonce submission (or last nonce if nonce is 0)
+ */
+extern "C" uint32_t hashlog_already_submittted(char* jobid, uint32_t nonce)
+{
+	uint32_t ret = 0;
+	uint64_t njobid = hextouint(jobid);
+	uint64_t key = (njobid << 32) + nonce;
+	if (nonce == 0) {
+		// search last submitted nonce for job
+		ret = hashlog_get_last_sent(jobid);
+	} else if (tlastshares.find(key) != tlastshares.end()) {
+		hashlog_data data = tlastshares[key];
+		ret = data.tm_sent;
+	}
+	return ret;
+}
+/**
+ * Store submitted nonces of a job
+ */
+extern "C" void hashlog_remember_submit(char* jobid, uint32_t nonce)
+{
+	uint64_t njobid = hextouint(jobid);
+	uint64_t keyall = (njobid << 32);
+	uint64_t key = keyall + nonce;
+	struct hashlog_data data;
+
+	data = tlastshares[keyall];
+	data.tm_upd = data.tm_sent = (uint32_t) time(NULL);
+	if (data.tm_add == 0)
+		data.tm_add = data.tm_upd;
+	tlastshares[key] = data;
+}
+
+/**
+ * Update job scanned range
+ */
+extern "C" void hashlog_remember_scan_range(char* jobid, uint32_t scanned_from, uint32_t scanned_to)
+{
+	uint64_t njobid = hextouint(jobid);
+	uint64_t keyall = (njobid << 32);
+	uint64_t range = hashlog_get_scan_range(jobid);
+	struct hashlog_data data;
+
+	// global scan range of a job
+	data = tlastshares[keyall];
+	if (range == 0) {
+		memset(&data, 0, sizeof(data));
+	} else {
+		// get min and max from all sent records
+		data.scanned_from = LO_DWORD(range);
+		data.scanned_to   = HI_DWORD(range);
+	}
+
+	if (data.tm_add == 0)
+		data.tm_add = (uint32_t) time(NULL);
+
+	data.last_from = scanned_from;
+
+	if (scanned_from < scanned_to) {
+		if (data.scanned_from == 0)
+			data.scanned_from = scanned_from ? scanned_from : 1; // min 1
+		else if (scanned_from < data.scanned_from) // || scanned_to == (data.scanned_from - 1)
+			data.scanned_from = scanned_from;
+		if (data.scanned_to == 0 || scanned_from == data.scanned_to + 1)
+			data.scanned_to = scanned_to;
+	}
+
+	data.tm_upd = (uint32_t) time(NULL);
+
+	tlastshares[keyall] = data;
+/* 	applog(LOG_BLUE, "job %s range : %x %x -> %x %x", jobid,
+		scanned_from, scanned_to, data.scanned_from, data.scanned_to); */
+}
+
+/**
+ * Returns the range of a job
+ * @return uint64_t to|from
+ */
+extern "C" uint64_t hashlog_get_scan_range(char* jobid)
+{
+	uint64_t ret = 0;
+	uint64_t njobid = hextouint(jobid);
+	uint64_t keypfx = (njobid << 32);
+	struct hashlog_data data;
+	data.scanned_from = 0;
+	data.scanned_to = 0;
+	std::map<uint64_t, hashlog_data>::iterator i = tlastshares.begin();
+	while (i != tlastshares.end()) {
+		if ((keypfx & i->first) == keypfx && i->second.scanned_to > ret) {
+			if (i->second.scanned_to > data.scanned_to)
+				data.scanned_to = i->second.scanned_to;
+			if (i->second.scanned_from < data.scanned_from || data.scanned_from == 0)
+				data.scanned_from = i->second.scanned_from;
+		}
+		i++;
+	}
+	ret = data.scanned_from;
+	ret += MK_HI64(data.scanned_to);
+	return ret;
+}
+
+/**
+ * Search last submitted nonce for a job
+ * @return max nonce
+ */
+extern "C" uint32_t hashlog_get_last_sent(char* jobid)
+{
+	uint32_t nonce = 0;
+	uint64_t njobid = hextouint(jobid);
+	uint64_t keypfx = (njobid << 32);
+	std::map<uint64_t, hashlog_data>::iterator i = tlastshares.begin();
+	while (i != tlastshares.end()) {
+		if ((keypfx & i->first) == keypfx && i->second.tm_sent > 0) {
+			nonce = LO_DWORD(i->first);
+		}
+		i++;
+	}
+	return nonce;
+}
+
+/**
+ * Remove entries of a job...
+ */
+extern "C" void hashlog_purge_job(char* jobid)
+{
+	int deleted = 0;
+	uint64_t njobid = hextouint(jobid);
+	uint64_t keypfx = (njobid << 32);
+	uint32_t sz = tlastshares.size();
+	std::map<uint64_t, hashlog_data>::iterator i = tlastshares.begin();
+	while (i != tlastshares.end()) {
+		if ((keypfx & i->first) == keypfx) {
+			deleted++;
+			tlastshares.erase(i++);
+		}
+		else ++i;
+	}
+	if (opt_debug && deleted) {
+		applog(LOG_DEBUG, "hashlog: purge job %s, del %d/%d", jobid, deleted, sz);
+	}
+}
+
+/**
+ * Remove old entries to reduce memory usage
+ */
+extern "C" void hashlog_purge_old(void)
+{
+	int deleted = 0;
+	uint32_t now = (uint32_t) time(NULL);
+	uint32_t sz = tlastshares.size();
+	std::map<uint64_t, hashlog_data>::iterator i = tlastshares.begin();
+	while (i != tlastshares.end()) {
+		if ((now - i->second.tm_sent) > LOG_PURGE_TIMEOUT) {
+			deleted++;
+			tlastshares.erase(i++);
+		}
+		else ++i;
+	}
+	if (opt_debug && deleted) {
+		applog(LOG_DEBUG, "hashlog: %d/%d purged", deleted, sz);
+	}
+}
+
+/**
+ * Reset the submitted nonces cache
+ */
+extern "C" void hashlog_purge_all(void)
+{
+	tlastshares.clear();
+}
+
+/**
+ * Used to debug ranges...
+ */
+extern "C" void hashlog_dump_job(char* jobid)
+{
+	if (opt_debug) {
+		int deleted = 0;
+		uint64_t njobid = hextouint(jobid);
+		uint64_t keypfx = (njobid << 32);
+		uint32_t sz = tlastshares.size();
+		std::map<uint64_t, hashlog_data>::iterator i = tlastshares.begin();
+		while (i != tlastshares.end()) {
+			if ((keypfx & i->first) == keypfx) {
+				applog(LOG_BLUE, "job %s range : %x %x %s added %x upd %x", jobid,
+					i->second.scanned_from, i->second.scanned_to,
+					i->second.tm_sent ? "sent" : "",
+					i->second.tm_add, i->second.tm_upd);/* */
+			}
+			i++;
+		}
+	}
+}

miner.h

@@ -241,6 +241,10 @@ extern int scanhash_fresh(int thr_id, uint32_t *pdata,
 	const uint32_t *ptarget, uint32_t max_nonce,
 	unsigned long *hashes_done);
 
+extern int scanhash_blake32(int thr_id, uint32_t *pdata,
+	const uint32_t *ptarget, uint32_t max_nonce,
+	unsigned long *hashes_done);
+
 extern int scanhash_nist5(int thr_id, uint32_t *pdata,
 	const uint32_t *ptarget, uint32_t max_nonce,
 	unsigned long *hashes_done);
@@ -281,6 +285,8 @@ struct work_restart {
 };
 
 extern bool opt_debug;
+extern bool opt_debug_rpc;
+extern bool opt_quiet;
 extern bool opt_protocol;
 extern int opt_timeout;
 extern bool want_longpoll;
@@ -311,7 +317,7 @@ extern uint16_t opt_vote;
 #define CL_BLK  "\x1B[22;30m" /* black */
 #define CL_RD2  "\x1B[22;31m" /* red */
 #define CL_GR2  "\x1B[22;32m" /* green */
-#define CL_BRW  "\x1B[22;33m" /* brown */
+#define CL_YL2  "\x1B[22;33m" /* dark yellow */
 #define CL_BL2  "\x1B[22;34m" /* blue */
 #define CL_MA2  "\x1B[22;35m" /* magenta */
 #define CL_CY2  "\x1B[22;36m" /* cyan */
@@ -320,7 +326,7 @@ extern uint16_t opt_vote;
 #define CL_GRY  "\x1B[01;30m" /* dark gray */
 #define CL_LRD  "\x1B[01;31m" /* light red */
 #define CL_LGR  "\x1B[01;32m" /* light green */
-#define CL_YL2  "\x1B[01;33m" /* yellow */
+#define CL_LYL  "\x1B[01;33m" /* tooltips */
 #define CL_LBL  "\x1B[01;34m" /* light blue */
 #define CL_LMA  "\x1B[01;35m" /* light magenta */
 #define CL_LCY  "\x1B[01;36m" /* light cyan */
@@ -372,6 +378,9 @@ struct stratum_ctx {
 	size_t xnonce2_size;
 	struct stratum_job job;
 	pthread_mutex_t work_lock;
+
+	int srvtime_diff;
+	int bloc_height;
 };
 
 bool stratum_socket_full(struct stratum_ctx *sctx, int timeout);
@@ -383,6 +392,16 @@ bool stratum_subscribe(struct stratum_ctx *sctx);
 bool stratum_authorize(struct stratum_ctx *sctx, const char *user, const char *pass);
 bool stratum_handle_method(struct stratum_ctx *sctx, const char *s);
 
+void hashlog_remember_submit(char* jobid, uint32_t nounce);
+void hashlog_remember_scan_range(char* jobid, uint32_t scanned_from, uint32_t scanned_to);
+uint32_t hashlog_already_submittted(char* jobid, uint32_t nounce);
+uint32_t hashlog_get_last_sent(char* jobid);
+uint64_t hashlog_get_scan_range(char* jobid);
+void hashlog_purge_old(void);
+void hashlog_purge_job(char* jobid);
+void hashlog_purge_all(void);
+void hashlog_dump_job(char* jobid);
+
 struct thread_q;
 
 extern struct thread_q *tq_new(void);
@@ -392,16 +411,22 @@ extern void *tq_pop(struct thread_q *tq, const struct timespec *abstime);
 extern void tq_freeze(struct thread_q *tq);
 extern void tq_thaw(struct thread_q *tq);
 
+void proper_exit(int reason);
+
+size_t time2str(char* buf, time_t timer);
+char* atime2str(time_t timer);
 
 void applog_hash(unsigned char *hash);
 
 void print_hash_tests(void);
-unsigned int jackpothash(void *state, const void *input);
+void animehash(void *state, const void *input);
+void blake32hash(void *output, const void *input);
+void fresh_hash(void *state, const void *input);
 void fugue256_hash(unsigned char* output, const unsigned char* input, int len);
 void heavycoin_hash(unsigned char* output, const unsigned char* input, int len);
+unsigned int jackpothash(void *state, const void *input);
 void groestlhash(void *state, const void *input);
 void myriadhash(void *state, const void *input);
-void fresh_hash(void *state, const void *input);
 void nist5hash(void *state, const void *input);
 void quarkhash(void *state, const void *input);
 void wcoinhash(void *state, const void *input);

missing

@@ -1,367 +0,0 @@
-#! /bin/sh
-# Common stub for a few missing GNU programs while installing.
-
-scriptversion=2006-05-10.23
-
-# Copyright (C) 1996, 1997, 1999, 2000, 2002, 2003, 2004, 2005, 2006
-#   Free Software Foundation, Inc.
-# Originally by Fran,cois Pinard <pinard@iro.umontreal.ca>, 1996.
-
-# 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 2, or (at your option)
-# any later version.
-
-# This program is distributed in the hope that it will be useful,
-# but WITHOUT ANY WARRANTY; without even the implied warranty of
-# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-# GNU General Public License for more details.
-
-# You should have received a copy of the GNU General Public License
-# along with this program; if not, write to the Free Software
-# Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA
-# 02110-1301, USA.
-
-# As a special exception to the GNU General Public License, if you
-# distribute this file as part of a program that contains a
-# configuration script generated by Autoconf, you may include it under
-# the same distribution terms that you use for the rest of that program.
-
-if test $# -eq 0; then
-  echo 1>&2 "Try \`$0 --help' for more information"
-  exit 1
-fi
-
-run=:
-sed_output='s/.* --output[ =]\([^ ]*\).*/\1/p'
-sed_minuso='s/.* -o \([^ ]*\).*/\1/p'
-
-# In the cases where this matters, `missing' is being run in the
-# srcdir already.
-if test -f configure.ac; then
-  configure_ac=configure.ac
-else
-  configure_ac=configure.in
-fi
-
-msg="missing on your system"
-
-case $1 in
---run)
-  # Try to run requested program, and just exit if it succeeds.
-  run=
-  shift
-  "$@" && exit 0
-  # Exit code 63 means version mismatch.  This often happens
-  # when the user try to use an ancient version of a tool on
-  # a file that requires a minimum version.  In this case we
-  # we should proceed has if the program had been absent, or
-  # if --run hadn't been passed.
-  if test $? = 63; then
-    run=:
-    msg="probably too old"
-  fi
-  ;;
-
-  -h|--h|--he|--hel|--help)
-    echo "\
-$0 [OPTION]... PROGRAM [ARGUMENT]...
-
-Handle \`PROGRAM [ARGUMENT]...' for when PROGRAM is missing, or return an
-error status if there is no known handling for PROGRAM.
-
-Options:
-  -h, --help      display this help and exit
-  -v, --version   output version information and exit
-  --run           try to run the given command, and emulate it if it fails
-
-Supported PROGRAM values:
-  aclocal      touch file \`aclocal.m4'
-  autoconf     touch file \`configure'
-  autoheader   touch file \`config.h.in'
-  autom4te     touch the output file, or create a stub one
-  automake     touch all \`Makefile.in' files
-  bison        create \`y.tab.[ch]', if possible, from existing .[ch]
-  flex         create \`lex.yy.c', if possible, from existing .c
-  help2man     touch the output file
-  lex          create \`lex.yy.c', if possible, from existing .c
-  makeinfo     touch the output file
-  tar          try tar, gnutar, gtar, then tar without non-portable flags
-  yacc         create \`y.tab.[ch]', if possible, from existing .[ch]
-
-Send bug reports to <bug-automake@gnu.org>."
-    exit $?
-    ;;
-
-  -v|--v|--ve|--ver|--vers|--versi|--versio|--version)
-    echo "missing $scriptversion (GNU Automake)"
-    exit $?
-    ;;
-
-  -*)
-    echo 1>&2 "$0: Unknown \`$1' option"
-    echo 1>&2 "Try \`$0 --help' for more information"
-    exit 1
-    ;;
-
-esac
-
-# Now exit if we have it, but it failed.  Also exit now if we
-# don't have it and --version was passed (most likely to detect
-# the program).
-case $1 in
-  lex|yacc)
-    # Not GNU programs, they don't have --version.
-    ;;
-
-  tar)
-    if test -n "$run"; then
-       echo 1>&2 "ERROR: \`tar' requires --run"
-       exit 1
-    elif test "x$2" = "x--version" || test "x$2" = "x--help"; then
-       exit 1
-    fi
-    ;;
-
-  *)
-    if test -z "$run" && ($1 --version) > /dev/null 2>&1; then
-       # We have it, but it failed.
-       exit 1
-    elif test "x$2" = "x--version" || test "x$2" = "x--help"; then
-       # Could not run --version or --help.  This is probably someone
-       # running `$TOOL --version' or `$TOOL --help' to check whether
-       # $TOOL exists and not knowing $TOOL uses missing.
-       exit 1
-    fi
-    ;;
-esac
-
-# If it does not exist, or fails to run (possibly an outdated version),
-# try to emulate it.
-case $1 in
-  aclocal*)
-    echo 1>&2 "\
-WARNING: \`$1' is $msg.  You should only need it if
-         you modified \`acinclude.m4' or \`${configure_ac}'.  You might want
-         to install the \`Automake' and \`Perl' packages.  Grab them from
-         any GNU archive site."
-    touch aclocal.m4
-    ;;
-
-  autoconf)
-    echo 1>&2 "\
-WARNING: \`$1' is $msg.  You should only need it if
-         you modified \`${configure_ac}'.  You might want to install the
-         \`Autoconf' and \`GNU m4' packages.  Grab them from any GNU
-         archive site."
-    touch configure
-    ;;
-
-  autoheader)
-    echo 1>&2 "\
-WARNING: \`$1' is $msg.  You should only need it if
-         you modified \`acconfig.h' or \`${configure_ac}'.  You might want
-         to install the \`Autoconf' and \`GNU m4' packages.  Grab them
-         from any GNU archive site."
-    files=`sed -n 's/^[ ]*A[CM]_CONFIG_HEADER(\([^)]*\)).*/\1/p' ${configure_ac}`
-    test -z "$files" && files="config.h"
-    touch_files=
-    for f in $files; do
-      case $f in
-      *:*) touch_files="$touch_files "`echo "$f" |
-				       sed -e 's/^[^:]*://' -e 's/:.*//'`;;
-      *) touch_files="$touch_files $f.in";;
-      esac
-    done
-    touch $touch_files
-    ;;
-
-  automake*)
-    echo 1>&2 "\
-WARNING: \`$1' is $msg.  You should only need it if
-         you modified \`Makefile.am', \`acinclude.m4' or \`${configure_ac}'.
-         You might want to install the \`Automake' and \`Perl' packages.
-         Grab them from any GNU archive site."
-    find . -type f -name Makefile.am -print |
-	   sed 's/\.am$/.in/' |
-	   while read f; do touch "$f"; done
-    ;;
-
-  autom4te)
-    echo 1>&2 "\
-WARNING: \`$1' is needed, but is $msg.
-         You might have modified some files without having the
-         proper tools for further handling them.
-         You can get \`$1' as part of \`Autoconf' from any GNU
-         archive site."
-
-    file=`echo "$*" | sed -n "$sed_output"`
-    test -z "$file" && file=`echo "$*" | sed -n "$sed_minuso"`
-    if test -f "$file"; then
-	touch $file
-    else
-	test -z "$file" || exec >$file
-	echo "#! /bin/sh"
-	echo "# Created by GNU Automake missing as a replacement of"
-	echo "#  $ $@"
-	echo "exit 0"
-	chmod +x $file
-	exit 1
-    fi
-    ;;
-
-  bison|yacc)
-    echo 1>&2 "\
-WARNING: \`$1' $msg.  You should only need it if
-         you modified a \`.y' file.  You may need the \`Bison' package
-         in order for those modifications to take effect.  You can get
-         \`Bison' from any GNU archive site."
-    rm -f y.tab.c y.tab.h
-    if test $# -ne 1; then
-        eval LASTARG="\${$#}"
-	case $LASTARG in
-	*.y)
-	    SRCFILE=`echo "$LASTARG" | sed 's/y$/c/'`
-	    if test -f "$SRCFILE"; then
-	         cp "$SRCFILE" y.tab.c
-	    fi
-	    SRCFILE=`echo "$LASTARG" | sed 's/y$/h/'`
-	    if test -f "$SRCFILE"; then
-	         cp "$SRCFILE" y.tab.h
-	    fi
-	  ;;
-	esac
-    fi
-    if test ! -f y.tab.h; then
-	echo >y.tab.h
-    fi
-    if test ! -f y.tab.c; then
-	echo 'main() { return 0; }' >y.tab.c
-    fi
-    ;;
-
-  lex|flex)
-    echo 1>&2 "\
-WARNING: \`$1' is $msg.  You should only need it if
-         you modified a \`.l' file.  You may need the \`Flex' package
-         in order for those modifications to take effect.  You can get
-         \`Flex' from any GNU archive site."
-    rm -f lex.yy.c
-    if test $# -ne 1; then
-        eval LASTARG="\${$#}"
-	case $LASTARG in
-	*.l)
-	    SRCFILE=`echo "$LASTARG" | sed 's/l$/c/'`
-	    if test -f "$SRCFILE"; then
-	         cp "$SRCFILE" lex.yy.c
-	    fi
-	  ;;
-	esac
-    fi
-    if test ! -f lex.yy.c; then
-	echo 'main() { return 0; }' >lex.yy.c
-    fi
-    ;;
-
-  help2man)
-    echo 1>&2 "\
-WARNING: \`$1' is $msg.  You should only need it if
-	 you modified a dependency of a manual page.  You may need the
-	 \`Help2man' package in order for those modifications to take
-	 effect.  You can get \`Help2man' from any GNU archive site."
-
-    file=`echo "$*" | sed -n "$sed_output"`
-    test -z "$file" && file=`echo "$*" | sed -n "$sed_minuso"`
-    if test -f "$file"; then
-	touch $file
-    else
-	test -z "$file" || exec >$file
-	echo ".ab help2man is required to generate this page"
-	exit 1
-    fi
-    ;;
-
-  makeinfo)
-    echo 1>&2 "\
-WARNING: \`$1' is $msg.  You should only need it if
-         you modified a \`.texi' or \`.texinfo' file, or any other file
-         indirectly affecting the aspect of the manual.  The spurious
-         call might also be the consequence of using a buggy \`make' (AIX,
-         DU, IRIX).  You might want to install the \`Texinfo' package or
-         the \`GNU make' package.  Grab either from any GNU archive site."
-    # The file to touch is that specified with -o ...
-    file=`echo "$*" | sed -n "$sed_output"`
-    test -z "$file" && file=`echo "$*" | sed -n "$sed_minuso"`
-    if test -z "$file"; then
-      # ... or it is the one specified with @setfilename ...
-      infile=`echo "$*" | sed 's/.* \([^ ]*\) *$/\1/'`
-      file=`sed -n '
-	/^@setfilename/{
-	  s/.* \([^ ]*\) *$/\1/
-	  p
-	  q
-	}' $infile`
-      # ... or it is derived from the source name (dir/f.texi becomes f.info)
-      test -z "$file" && file=`echo "$infile" | sed 's,.*/,,;s,.[^.]*$,,'`.info
-    fi
-    # If the file does not exist, the user really needs makeinfo;
-    # let's fail without touching anything.
-    test -f $file || exit 1
-    touch $file
-    ;;
-
-  tar)
-    shift
-
-    # We have already tried tar in the generic part.
-    # Look for gnutar/gtar before invocation to avoid ugly error
-    # messages.
-    if (gnutar --version > /dev/null 2>&1); then
-       gnutar "$@" && exit 0
-    fi
-    if (gtar --version > /dev/null 2>&1); then
-       gtar "$@" && exit 0
-    fi
-    firstarg="$1"
-    if shift; then
-	case $firstarg in
-	*o*)
-	    firstarg=`echo "$firstarg" | sed s/o//`
-	    tar "$firstarg" "$@" && exit 0
-	    ;;
-	esac
-	case $firstarg in
-	*h*)
-	    firstarg=`echo "$firstarg" | sed s/h//`
-	    tar "$firstarg" "$@" && exit 0
-	    ;;
-	esac
-    fi
-
-    echo 1>&2 "\
-WARNING: I can't seem to be able to run \`tar' with the given arguments.
-         You may want to install GNU tar or Free paxutils, or check the
-         command line arguments."
-    exit 1
-    ;;
-
-  *)
-    echo 1>&2 "\
-WARNING: \`$1' is needed, and is $msg.
-         You might have modified some files without having the
-         proper tools for further handling them.  Check the \`README' file,
-         it often tells you about the needed prerequisites for installing
-         this package.  You may also peek at any GNU archive site, in case
-         some other package would contain this missing \`$1' program."
-    exit 1
-    ;;
-esac
-
-exit 0
-
-# Local variables:
-# eval: (add-hook 'write-file-hooks 'time-stamp)
-# time-stamp-start: "scriptversion="
-# time-stamp-format: "%:y-%02m-%02d.%02H"
-# time-stamp-end: "$"
-# End:

quark/animecoin.cu

@@ -57,7 +57,7 @@ extern void quark_compactTest_single_false_cpu_hash_64(int thr_id, int threads,
 											int order);
 
 // Original Quarkhash Funktion aus einem miner Quelltext
-inline void animehash(void *state, const void *input)
+extern "C" void animehash(void *state, const void *input)
 {
     sph_blake512_context ctx_blake;
     sph_bmw512_context ctx_bmw;

quark/cuda_checkhash.cu

@@ -6,8 +6,8 @@
 // Hash Target gegen das wir testen sollen
 __constant__ uint32_t pTarget[8];
 
-uint32_t *d_resNounce[8];
+static uint32_t *d_resNounce[8];
-uint32_t *h_resNounce[8];
+static uint32_t *h_resNounce[8];
 
 // aus heavy.cu
 extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);

util.c

@@ -115,7 +115,7 @@ void applog(int prio, const char *fmt, ...)
 			case LOG_WARNING: color = CL_YLW; break;
 			case LOG_NOTICE:  color = CL_WHT; break;
 			case LOG_INFO:    color = ""; break;
-			case LOG_DEBUG:   color = ""; break;
+			case LOG_DEBUG:   color = CL_GRY; break;
 
 			case LOG_BLUE:
 				prio = LOG_NOTICE;
@@ -559,7 +559,7 @@ bool fulltest(const uint32_t *hash, const uint32_t *target)
 		}
 	}
 
-	if (opt_debug) {
+	if (!rc && opt_debug) {
 		uint32_t hash_be[8], target_be[8];
 		char *hash_str, *target_str;
 		
@@ -572,7 +572,7 @@ bool fulltest(const uint32_t *hash, const uint32_t *target)
 
 		applog(LOG_DEBUG, "DEBUG: %s\nHash:   %s\nTarget: %s",
 			rc ? "hash <= target"
-			   : "hash > target (false positive)",
+			   : CL_YLW "hash > target (false positive)" CL_N,
 			hash_str,
 			target_str);
 
@@ -1011,12 +1011,13 @@ out:
 
 static bool stratum_notify(struct stratum_ctx *sctx, json_t *params)
 {
-	const char *job_id, *prevhash, *coinb1, *coinb2, *version, *nbits, *ntime, *nreward;
+	const char *job_id, *prevhash, *coinb1, *coinb2, *version, *nbits, *stime, *nreward;
 	size_t coinb1_size, coinb2_size;
 	bool clean, ret = false;
 	int merkle_count, i;
 	json_t *merkle_arr;
 	unsigned char **merkle;
+	int ntime;
 
 	job_id = json_string_value(json_array_get(params, 0));
 	prevhash = json_string_value(json_array_get(params, 1));
@@ -1028,16 +1029,26 @@ static bool stratum_notify(struct stratum_ctx *sctx, json_t *params)
 	merkle_count = json_array_size(merkle_arr);
 	version = json_string_value(json_array_get(params, 5));
 	nbits = json_string_value(json_array_get(params, 6));
-	ntime = json_string_value(json_array_get(params, 7));
+	stime = json_string_value(json_array_get(params, 7));
 	clean = json_is_true(json_array_get(params, 8));
 	nreward = json_string_value(json_array_get(params, 9));
 
-	if (!job_id || !prevhash || !coinb1 || !coinb2 || !version || !nbits || !ntime ||
+	if (!job_id || !prevhash || !coinb1 || !coinb2 || !version || !nbits || !stime ||
 	    strlen(prevhash) != 64 || strlen(version) != 8 ||
-	    strlen(nbits) != 8 || strlen(ntime) != 8) {
+	    strlen(nbits) != 8 || strlen(stime) != 8) {
 		applog(LOG_ERR, "Stratum notify: invalid parameters");
 		goto out;
 	}
+
+	/* store stratum server time diff */
+	hex2bin((unsigned char *)&ntime, stime, 4);
+	ntime = swab32(ntime) - time(0);
+	if (ntime > sctx->srvtime_diff) {
+		sctx->srvtime_diff = ntime;
+		if (!opt_quiet)
+			applog(LOG_DEBUG, "stratum time is at least %ds in the future", ntime);
+	}
+
 	merkle = (unsigned char**)malloc(merkle_count * sizeof(char *));
 	for (i = 0; i < merkle_count; i++) {
 		const char *s = json_string_value(json_array_get(merkle_arr, i));
@@ -1058,10 +1069,13 @@ static bool stratum_notify(struct stratum_ctx *sctx, json_t *params)
 	coinb2_size = strlen(coinb2) / 2;
 	sctx->job.coinbase_size = coinb1_size + sctx->xnonce1_size +
 	                          sctx->xnonce2_size + coinb2_size;
+
 	sctx->job.coinbase = (unsigned char*)realloc(sctx->job.coinbase, sctx->job.coinbase_size);
 	sctx->job.xnonce2 = sctx->job.coinbase + coinb1_size + sctx->xnonce1_size;
 	hex2bin(sctx->job.coinbase, coinb1, coinb1_size);
 	memcpy(sctx->job.coinbase + coinb1_size, sctx->xnonce1, sctx->xnonce1_size);
+
+	sctx->bloc_height = le16dec((uint8_t*) sctx->job.coinbase + 43);
 	if (!sctx->job.job_id || strcmp(sctx->job.job_id, job_id))
 		memset(sctx->job.xnonce2, 0, sctx->xnonce2_size);
 	hex2bin(sctx->job.xnonce2 + sctx->xnonce2_size, coinb2, coinb2_size);
@@ -1078,7 +1092,7 @@ static bool stratum_notify(struct stratum_ctx *sctx, json_t *params)
 
 	hex2bin(sctx->job.version, version, 4);
 	hex2bin(sctx->job.nbits, nbits, 4);
-	hex2bin(sctx->job.ntime, ntime, 4);
+	hex2bin(sctx->job.ntime, stime, 4);
 	if(nreward != NULL)
 	{
 		if(strlen(nreward) == 4)
@@ -1205,6 +1219,10 @@ bool stratum_handle_method(struct stratum_ctx *sctx, const char *s)
 	id = json_object_get(val, "id");
 	params = json_object_get(val, "params");
 
+	if (opt_debug_rpc) {
+		applog(LOG_DEBUG, "method: %s", s);
+	}
+
 	if (!strcasecmp(method, "mining.notify")) {
 		ret = stratum_notify(sctx, params);
 		goto out;
@@ -1346,6 +1364,29 @@ out:
 	return rval;
 }
 
+/**
+ * @param buf char[9] mini
+ * @param time_t timer to convert
+ */
+size_t time2str(char* buf, time_t timer)
+{
+	struct tm* tm_info;
+	tm_info = localtime(&timer);
+	return strftime(buf, 19, "%H:%M:%S", tm_info);
+}
+
+/**
+ * Alloc and returns time string (to be freed)
+ * @param time_t timer to convert
+ */
+char* atime2str(time_t timer)
+{
+	char* buf = (char*) malloc(16);
+	memset(buf, 0, 16);
+	time2str(buf, timer);
+	return buf;
+}
+
 /* sprintf can be used in applog */
 static char* format_hash(char* buf, unsigned char *hash)
 {
@@ -1368,11 +1409,24 @@ extern void applog_hash(unsigned char *hash)
 
 void print_hash_tests(void)
 {
-	unsigned char buf[128], hash[128], s[128];
+	char s[128] = {'\0'};
+	unsigned char buf[128], hash[128];
 	memset(buf, 0, sizeof buf);
 
 	printf(CL_WHT "CPU HASH ON EMPTY BUFFER RESULTS:" CL_N "\n");
 
+	memset(hash, 0, sizeof hash);
+	animehash(&hash[0], &buf[0]);
+	printpfx("anime", hash);
+
+	memset(hash, 0, sizeof hash);
+	blake32hash(&hash[0], &buf[0]);
+	printpfx("blake", hash);
+
+	memset(hash, 0, sizeof hash);
+	fresh_hash(&hash[0], &buf[0]);
+	printpfx("fresh", hash);
+
 	memset(hash, 0, sizeof hash);
 	fugue256_hash(&hash[0], &buf[0], 32);
 	printpfx("fugue256", hash);
@@ -1401,10 +1455,6 @@ void print_hash_tests(void)
 	quarkhash(&hash[0], &buf[0]);
 	printpfx("quark", hash);
 
-	memset(hash, 0, sizeof hash);
-	fresh_hash(&hash[0], &buf[0]);
-	printpfx("fresh", hash);
-
 	memset(hash, 0, sizeof hash);
 	wcoinhash(&hash[0], &buf[0]);
 	printpfx("whirl", hash);