From 3e44e5230076413d2a840d1dc67450b8cf3e3541 Mon Sep 17 00:00:00 2001 From: Tanguy Pruvot Date: Wed, 3 Sep 2014 14:20:49 +0200 Subject: [PATCH] Release v1.4 with blake (NEOS) Blake256: squashed commit... Squashed commit of the following: commit c370208bc92ef16557f66e5391faf2b1ad47726f Author: Tanguy Pruvot Date: Wed Sep 3 13:53:01 2014 +0200 hashlog: prepare store of scanned range commit e2cf49a5e956f03deafd266d1a0dd087a2041c99 Author: Tanguy Pruvot Date: Wed Sep 3 12:54:13 2014 +0200 stratum: store server time offset in context commit 1a4391d7ff21397a128abf031f92733a8ac47437 Author: Tanguy Pruvot Date: Tue Sep 2 12:40:52 2014 +0200 hashlog: prevent double computing on jobs already done commit 049e57730116685755bd3ff214f0793cce7c773b Author: Tanguy Pruvot Date: Wed Sep 3 09:49:14 2014 +0200 tmp blake log commit 43d3e93e1a97e569ead2437f759c6b8423d30c0a Author: Tanguy Pruvot Date: Wed Sep 3 09:29:51 2014 +0200 blake: set a max throughput commit 7e595a36ea69027c8a28023399540a761e7686c3 Author: Tanguy Pruvot 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 de80c7e9d1448f15541d08c5dbbf372d5bfeba48 Author: Tanguy Pruvot 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 2d42ae6de586a6ae8cbfd01806a273fd5cc4b262 Author: Tanguy Pruvot 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 commit 1b8c3c12fa5bb83afbb02f9d5f60586939f36d86 Author: Tanguy Pruvot Date: Tue Sep 2 03:38:57 2014 +0200 debug: a new boolean to log or not json rpc data commit 1f99aae0ff621f4f85f119d811a3f1a8d2204f60 Author: Tanguy Pruvot 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 530732458add6c4c3836606d028930f3581c0a5f Author: Tanguy Pruvot 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 0aeac878ef60840f3123354037cd56a89d2e94e6 Author: Tanguy Pruvot 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 4a52d0553b0076b984be480725fa67689c544647 Author: Tanguy Pruvot Date: Mon Sep 1 10:22:32 2014 +0200 debug: show json methods, hide hash/target if ok commit 1fb9becc1f2b6a15d8ccea4d8314df9ddf0af4ed Author: Tanguy Pruvot Date: Mon Sep 1 08:44:19 2014 +0200 cpu-miner: sort algos by name, show reject reason commit bfe96c49b0bf321ed0776cb1cf31c4fe8a0a8b8d Author: Tanguy Pruvot Date: Mon Aug 25 11:21:06 2014 +0200 release 1.4, update README... commit c17d11e37758c37762a7664a731fda6e9a5454b1 Author: Tanguy Pruvot 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 | 23 ++- README.md | 31 +++- README.txt | 15 ++ blake32.cu | 351 ++++++++++++++++++++++++++++++++++++++ ccminer.vcxproj | 8 +- ccminer.vcxproj.filters | 8 +- configure.ac | 6 +- configure.sh | 10 +- cpu-miner.c | 300 ++++++++++++++++++++++++-------- cpuminer-config.h | 6 +- hashlog.cpp | 226 +++++++++++++++++++++++++ miner.h | 33 +++- missing | 367 ---------------------------------------- quark/animecoin.cu | 2 +- quark/cuda_checkhash.cu | 4 +- util.c | 76 +++++++-- 16 files changed, 984 insertions(+), 482 deletions(-) create mode 100644 blake32.cu create mode 100644 hashlog.cpp delete mode 100644 missing diff --git a/Makefile.am b/Makefile.am index 5e539a9490..875f8b19ce 100644 --- a/Makefile.am +++ b/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 $< diff --git a/README.md b/README.md index af8fd12eb2..d3836eb037 100644 --- a/README.md +++ b/README.md @@ -3,5 +3,32 @@ ccminer Christian Buchner's & Christian H.'s CUDA miner project -Fork by tpruvot@github with X14 support - BTC donation address: 1AJdfCpLWPNoAMDfHF1wD5y8VgKSSTHxPo \ No newline at end of file +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. + diff --git a/README.txt b/README.txt index 32b2599300..e6fe248d87 100644 --- a/README.txt +++ b/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 diff --git a/blake32.cu b/blake32.cu new file mode 100644 index 0000000000..8be5205684 --- /dev/null +++ b/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 +#include +} + +/* 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<<>>(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; +} diff --git a/ccminer.vcxproj b/ccminer.vcxproj index 8bada54f22..7590d94c2a 100644 --- a/ccminer.vcxproj +++ b/ccminer.vcxproj @@ -243,6 +243,7 @@ copy "$(CudaToolkitBinDir)\cudart64*.dll" "$(OutDir)" + @@ -397,6 +398,11 @@ copy "$(CudaToolkitBinDir)\cudart64*.dll" "$(OutDir)" %(AdditionalOptions) 64 + + 64 + --ptxas-options=-O2 %(AdditionalOptions) + %(AdditionalOptions) + --ptxas-options=-O2 %(AdditionalOptions) %(AdditionalOptions) @@ -556,4 +562,4 @@ copy "$(CudaToolkitBinDir)\cudart64*.dll" "$(OutDir)" - \ No newline at end of file + diff --git a/ccminer.vcxproj.filters b/ccminer.vcxproj.filters index c972707513..93e331c814 100644 --- a/ccminer.vcxproj.filters +++ b/ccminer.vcxproj.filters @@ -180,6 +180,9 @@ Source Files + + Source Files + @@ -436,5 +439,8 @@ Source Files\CUDA\x17 + + Source Files\CUDA + - + \ No newline at end of file diff --git a/configure.ac b/configure.ac index 14e946863a..f7924d4a4d 100644 --- a/configure.ac +++ b/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" diff --git a/configure.sh b/configure.sh index 134abd163b..9c8b0218d7 100755 --- a/configure.sh +++ b/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 + diff --git a/cpu-miner.c b/cpu-miner.c index a55f051031..dc86a3e433 100644 --- a/cpu-miner.c +++ b/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_MJOLLNIR, /* Mjollnir hash */ + ALGO_ANIME, + 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", - "mjollnir", + "anime", + "blake", + "fresh", "fugue256", "groestl", - "myr-gr", + "heavy", "jackpot", - "quark", - "anime", - "fresh", + "mjollnir", + "myr-gr", "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\ + 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\ - groestl Groestlcoin hash\n\ myr-gr Myriad-Groestl hash\n\ - jackpot Jackpot hash\n\ - quark Quark hash\n\ - anime Animecoin hash\n\ - fresh Freshcoin hash (shavite 80)\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 @@ static bool work_decode(const json_t *val, struct work *work) 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) - applog(LOG_DEBUG, "DEBUG: reject reason: %s", reason); + if (reason && !opt_quiet) { + 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", - work->job_id, xnonce2str, swab32(work->data[17])); + applog(LOG_DEBUG, "DEBUG: job_id=%s xnonce2=%s time=%s", + 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) - sleep(1); + while (time(NULL) >= (g_work_time + opt_scantime) && !work_done) + 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 || - time(NULL) >= g_work_time + LP_SCANTIME*3/4 || - work.data[19] >= end_nonce)) { + if (time(NULL) - g_work_time >= min_scantime || + (*nonceptr) >= 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); - } - if (have_stratum) { - pthread_mutex_unlock(&g_work_lock); - continue; + g_work_time = time(NULL); } } - if (memcmp(work.data, g_work.data, 76)) { + if (memcmp(work.data, g_work.data, wcmplen)) { + if (opt_debug) { + applog(LOG_DEBUG, "job %s work updated", g_work.job_id); + for (int n=0; n end_nonce) + + if (max64 <= 0) { + 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 @@ static void *stratum_thread(void *userdata) 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(); - exit(0); + proper_exit(0); break; case SIGTERM: applog(LOG_INFO, "SIGTERM received, exiting"); - cuda_devicereset(); - exit(0); + proper_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(); - exit(0); + proper_exit(0); break; case CTRL_BREAK_EVENT: applog(LOG_INFO, "CTRL_BREAK_EVENT received, exiting"); - cuda_devicereset(); - exit(0); + proper_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; } diff --git a/cpuminer-config.h b/cpuminer-config.h index 0d0f042ae2..0fafa855c3 100644 --- a/cpuminer-config.h +++ b/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 */ diff --git a/hashlog.cpp b/hashlog.cpp new file mode 100644 index 0000000000..7a679e80b3 --- /dev/null +++ b/hashlog.cpp @@ -0,0 +1,226 @@ +#include +#include +#include + +#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 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::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::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::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::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::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++; + } + } +} \ No newline at end of file diff --git a/miner.h b/miner.h index f3d4299681..6ce4ca8b81 100644 --- a/miner.h +++ b/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); diff --git a/missing b/missing deleted file mode 100644 index 1c8ff7049d..0000000000 --- a/missing +++ /dev/null @@ -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 , 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 ." - 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: diff --git a/quark/animecoin.cu b/quark/animecoin.cu index c19275da5d..4b2d097840 100644 --- a/quark/animecoin.cu +++ b/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; diff --git a/quark/cuda_checkhash.cu b/quark/cuda_checkhash.cu index 3c41a02fc5..1ce25ecffc 100644 --- a/quark/cuda_checkhash.cu +++ b/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]; -uint32_t *h_resNounce[8]; +static uint32_t *d_resNounce[8]; +static uint32_t *h_resNounce[8]; // aus heavy.cu extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id); diff --git a/util.c b/util.c index 79a7a24032..f451d95000 100644 --- a/util.c +++ b/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 @@ bool stratum_authorize(struct stratum_ctx *sctx, const char *user, const char *p 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 @@ void *tq_pop(struct thread_q *tq, const struct timespec *abstime) 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);