diff --git a/Makefile.am b/Makefile.am
index 4ed9679..0ff48b3 100644
--- a/Makefile.am
+++ b/Makefile.am
@@ -64,7 +64,7 @@ ccminer_SOURCES = elist.h miner.h compat.h \
qubit/qubit.cu qubit/qubit_luffa512.cu qubit/deep.cu qubit/luffa.cu \
x11/x11.cu x11/fresh.cu x11/cuda_x11_luffa512.cu x11/cuda_x11_cubehash512.cu \
x11/cuda_x11_shavite512.cu x11/cuda_x11_simd512.cu x11/cuda_x11_echo.cu \
- x11/cuda_x11_luffa512_Cubehash.cu x11/x11evo.cu x11/timetravel.cu \
+ x11/cuda_x11_luffa512_Cubehash.cu x11/x11evo.cu x11/timetravel.cu x11/bitcore.cu \
x13/x13.cu x13/cuda_x13_hamsi512.cu x13/cuda_x13_fugue512.cu \
x15/x14.cu x15/x15.cu x15/cuda_x14_shabal512.cu x15/cuda_x15_whirlpool.cu \
x15/whirlpool.cu x15/cuda_x15_whirlpool_sm3.cu \
diff --git a/algos.h b/algos.h
index ea27f7c..926f2a4 100644
--- a/algos.h
+++ b/algos.h
@@ -45,6 +45,7 @@ enum sha_algos {
ALGO_SKEIN2,
ALGO_S3,
ALGO_TIMETRAVEL,
+ ALGO_BITCORE,
ALGO_X11EVO,
ALGO_X11,
ALGO_X13,
@@ -105,6 +106,7 @@ static const char *algo_names[] = {
"skein2",
"s3",
"timetravel",
+ "bitcore",
"x11evo",
"x11",
"x13",
@@ -159,6 +161,8 @@ static inline int algo_to_int(char* arg)
i = ALGO_SHA256D;
else if (!strcasecmp("thorsriddle", arg))
i = ALGO_VELTOR;
+ else if (!strcasecmp("timetravel10", arg))
+ i = ALGO_BITCORE;
else if (!strcasecmp("whirl", arg))
i = ALGO_WHIRLPOOL;
else if (!strcasecmp("ziftr", arg))
diff --git a/bench.cpp b/bench.cpp
index 5fa2791..12854a6 100644
--- a/bench.cpp
+++ b/bench.cpp
@@ -94,6 +94,7 @@ void algo_free_all(int thr_id)
free_scrypt(thr_id);
free_scrypt_jane(thr_id);
free_timetravel(thr_id);
+ free_bitcore(thr_id);
}
// benchmark all algos (called once per mining thread)
diff --git a/ccminer.cpp b/ccminer.cpp
index 04b166a..1efdfbc 100644
--- a/ccminer.cpp
+++ b/ccminer.cpp
@@ -224,6 +224,7 @@ Usage: " PROGRAM_NAME " [OPTIONS]\n\
Options:\n\
-a, --algo=ALGO specify the hash algorithm to use\n\
bastion Hefty bastion\n\
+ bitcore Timetravel-10\n\
blake Blake 256 (SFR)\n\
blake2s Blake2-S 256 (NEVA)\n\
blakecoin Fast Blake 256 (8 rounds)\n\
@@ -1625,6 +1626,7 @@ static bool stratum_gen_work(struct stratum_ctx *sctx, struct work *work)
case ALGO_LYRA2v2:
case ALGO_LYRA2Z:
case ALGO_TIMETRAVEL:
+ case ALGO_BITCORE:
work_set_target(work, sctx->job.diff / (256.0 * opt_difficulty));
break;
case ALGO_KECCAK:
@@ -2129,6 +2131,7 @@ static void *miner_thread(void *userdata)
case ALGO_LYRA2v2:
case ALGO_S3:
case ALGO_TIMETRAVEL:
+ case ALGO_BITCORE:
case ALGO_X11EVO:
case ALGO_X11:
case ALGO_X13:
@@ -2347,6 +2350,9 @@ static void *miner_thread(void *userdata)
case ALGO_TIMETRAVEL:
rc = scanhash_timetravel(thr_id, &work, max_nonce, &hashes_done);
break;
+ case ALGO_BITCORE:
+ rc = scanhash_bitcore(thr_id, &work, max_nonce, &hashes_done);
+ break;
case ALGO_X11EVO:
rc = scanhash_x11evo(thr_id, &work, max_nonce, &hashes_done);
break;
diff --git a/ccminer.vcxproj b/ccminer.vcxproj
index afddcc2..2b26598 100644
--- a/ccminer.vcxproj
+++ b/ccminer.vcxproj
@@ -538,6 +538,7 @@
+
diff --git a/ccminer.vcxproj.filters b/ccminer.vcxproj.filters
index 775da32..97302bb 100644
--- a/ccminer.vcxproj.filters
+++ b/ccminer.vcxproj.filters
@@ -739,6 +739,9 @@
Source Files\CUDA\x11
+
+ Source Files\CUDA\x11
+
Source Files\CUDA\x11
diff --git a/miner.h b/miner.h
index c4c739f..221d6ee 100644
--- a/miner.h
+++ b/miner.h
@@ -308,6 +308,7 @@ extern int scanhash_skeincoin(int thr_id, struct work* work, uint32_t max_nonce,
extern int scanhash_skein2(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done);
extern int scanhash_s3(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done);
extern int scanhash_timetravel(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done);
+extern int scanhash_bitcore(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done);
extern int scanhash_vanilla(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done, int8_t blake_rounds);
extern int scanhash_veltor(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done);
extern int scanhash_whirl(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done);
@@ -363,6 +364,7 @@ extern void free_skeincoin(int thr_id);
extern void free_skein2(int thr_id);
extern void free_s3(int thr_id);
extern void free_timetravel(int thr_id);
+extern void free_bitcore(int thr_id);
extern void free_vanilla(int thr_id);
extern void free_veltor(int thr_id);
extern void free_whirl(int thr_id);
@@ -887,6 +889,7 @@ void skeincoinhash(void *output, const void *input);
void skein2hash(void *output, const void *input);
void s3hash(void *output, const void *input);
void timetravel_hash(void *output, const void *input);
+void bitcore_hash(void *output, const void *input);
void veltorhash(void *output, const void *input);
void wcoinhash(void *state, const void *input);
void whirlxHash(void *state, const void *input);
diff --git a/util.cpp b/util.cpp
index a01bf30..1e945c2 100644
--- a/util.cpp
+++ b/util.cpp
@@ -2252,6 +2252,9 @@ void print_hash_tests(void)
timetravel_hash(&hash[0], &buf[0]);
printpfx("timetravel", hash);
+ bitcore_hash(&hash[0], &buf[0]);
+ printpfx("bitcore", hash);
+
blake256hash(&hash[0], &buf[0], 8);
printpfx("vanilla", hash);
diff --git a/x11/bitcore.cu b/x11/bitcore.cu
new file mode 100644
index 0000000..7873967
--- /dev/null
+++ b/x11/bitcore.cu
@@ -0,0 +1,439 @@
+/**
+ * Timetravel-10 (bitcore) CUDA implementation
+ * by tpruvot@github - May 2017
+ */
+
+#include
+#include
+#include
+
+#define HASH_FUNC_BASE_TIMESTAMP 1492973331U
+#define HASH_FUNC_COUNT 10
+#define HASH_FUNC_COUNT_PERMUTATIONS 40320U
+
+extern "C" {
+#include "sph/sph_blake.h"
+#include "sph/sph_bmw.h"
+#include "sph/sph_groestl.h"
+#include "sph/sph_skein.h"
+#include "sph/sph_jh.h"
+#include "sph/sph_keccak.h"
+#include "sph/sph_luffa.h"
+#include "sph/sph_cubehash.h"
+#include "sph/sph_shavite.h"
+#include "sph/sph_simd.h"
+#if HASH_FUNC_COUNT > 10
+#include "sph/sph_echo.h"
+#endif
+}
+
+#include "miner.h"
+#include "cuda_helper.h"
+#include "cuda_x11.h"
+
+static uint32_t *d_hash[MAX_GPUS];
+
+enum Algo {
+ BLAKE = 0,
+ BMW,
+ GROESTL,
+ SKEIN,
+ JH,
+ KECCAK,
+ LUFFA,
+ CUBEHASH,
+ SHAVITE,
+ SIMD,
+#if HASH_FUNC_COUNT > 10
+ ECHO,
+#endif
+ MAX_ALGOS_COUNT
+};
+
+inline void swap8(uint8_t *a, uint8_t *b)
+{
+ uint8_t t = *a;
+ *a = *b;
+ *b = t;
+}
+
+inline void initPerm(uint8_t n[], int count)
+{
+ for (int i = 0; i < count; i++)
+ n[i] = i;
+}
+
+static int nextPerm(uint8_t n[], int count)
+{
+ int tail, i, j;
+
+ if (count <= 1)
+ return 0;
+
+ for (i = count - 1; i>0 && n[i - 1] >= n[i]; i--);
+ tail = i;
+
+ if (tail > 0) {
+ for (j = count - 1; j>tail && n[j] <= n[tail - 1]; j--);
+ swap8(&n[tail - 1], &n[j]);
+ }
+
+ for (i = tail, j = count - 1; i= 10)
+ sprintf(sptr, "%c", 'A' + (algoList[j] - 10));
+ else
+ sprintf(sptr, "%u", (uint32_t) algoList[j]);
+ sptr++;
+ }
+ *sptr = '\0';
+}
+
+static __thread uint32_t s_ntime = 0;
+static uint32_t s_sequence = UINT32_MAX;
+static uint8_t s_firstalgo = 0xFF;
+static char hashOrder[HASH_FUNC_COUNT + 1] = { 0 };
+
+#define INITIAL_DATE HASH_FUNC_BASE_TIMESTAMP
+static inline uint32_t getCurrentAlgoSeq(uint32_t ntime)
+{
+ // unlike x11evo, the permutation changes often (with ntime)
+ return (uint32_t) (ntime - INITIAL_DATE) % HASH_FUNC_COUNT_PERMUTATIONS;
+}
+
+// To finish...
+static void get_travel_order(uint32_t ntime, char *permstr)
+{
+ uint32_t seq = getCurrentAlgoSeq(ntime);
+ if (s_sequence != seq) {
+ getAlgoString(permstr, seq);
+ s_sequence = seq;
+ }
+}
+
+// CPU Hash
+extern "C" void bitcore_hash(void *output, const void *input)
+{
+ uint32_t _ALIGN(64) hash[64/4] = { 0 };
+
+ sph_blake512_context ctx_blake;
+ sph_bmw512_context ctx_bmw;
+ sph_groestl512_context ctx_groestl;
+ sph_skein512_context ctx_skein;
+ sph_jh512_context ctx_jh;
+ sph_keccak512_context ctx_keccak;
+ sph_luffa512_context ctx_luffa1;
+ sph_cubehash512_context ctx_cubehash1;
+ sph_shavite512_context ctx_shavite1;
+ sph_simd512_context ctx_simd1;
+#if HASH_FUNC_COUNT > 10
+ sph_echo512_context ctx_echo1;
+#endif
+
+ if (s_sequence == UINT32_MAX) {
+ uint32_t *data = (uint32_t*) input;
+ const uint32_t ntime = (opt_benchmark || !data[17]) ? (uint32_t) time(NULL) : data[17];
+ get_travel_order(ntime, hashOrder);
+ }
+
+ void *in = (void*) input;
+ int size = 80;
+
+ const int hashes = (int) strlen(hashOrder);
+
+ for (int i = 0; i < hashes; i++)
+ {
+ const char elem = hashOrder[i];
+ uint8_t algo = elem >= 'A' ? elem - 'A' + 10 : elem - '0';
+
+ if (i > 0) {
+ in = (void*) hash;
+ size = 64;
+ }
+
+ switch (algo) {
+ case BLAKE:
+ sph_blake512_init(&ctx_blake);
+ sph_blake512(&ctx_blake, in, size);
+ sph_blake512_close(&ctx_blake, hash);
+ break;
+ case BMW:
+ sph_bmw512_init(&ctx_bmw);
+ sph_bmw512(&ctx_bmw, in, size);
+ sph_bmw512_close(&ctx_bmw, hash);
+ break;
+ case GROESTL:
+ sph_groestl512_init(&ctx_groestl);
+ sph_groestl512(&ctx_groestl, in, size);
+ sph_groestl512_close(&ctx_groestl, hash);
+ break;
+ case SKEIN:
+ sph_skein512_init(&ctx_skein);
+ sph_skein512(&ctx_skein, in, size);
+ sph_skein512_close(&ctx_skein, hash);
+ break;
+ case JH:
+ sph_jh512_init(&ctx_jh);
+ sph_jh512(&ctx_jh, in, size);
+ sph_jh512_close(&ctx_jh, hash);
+ break;
+ case KECCAK:
+ sph_keccak512_init(&ctx_keccak);
+ sph_keccak512(&ctx_keccak, in, size);
+ sph_keccak512_close(&ctx_keccak, hash);
+ break;
+ case LUFFA:
+ sph_luffa512_init(&ctx_luffa1);
+ sph_luffa512(&ctx_luffa1, in, size);
+ sph_luffa512_close(&ctx_luffa1, hash);
+ break;
+ case CUBEHASH:
+ sph_cubehash512_init(&ctx_cubehash1);
+ sph_cubehash512(&ctx_cubehash1, in, size);
+ sph_cubehash512_close(&ctx_cubehash1, hash);
+ break;
+ case SHAVITE:
+ sph_shavite512_init(&ctx_shavite1);
+ sph_shavite512(&ctx_shavite1, in, size);
+ sph_shavite512_close(&ctx_shavite1, hash);
+ break;
+ case SIMD:
+ sph_simd512_init(&ctx_simd1);
+ sph_simd512(&ctx_simd1, in, size);
+ sph_simd512_close(&ctx_simd1, hash);
+ break;
+#if HASH_FUNC_COUNT > 10
+ case ECHO:
+ sph_echo512_init(&ctx_echo1);
+ sph_echo512(&ctx_echo1, in, size);
+ sph_echo512_close(&ctx_echo1, hash);
+ break;
+#endif
+ }
+ }
+
+ memcpy(output, hash, 32);
+}
+
+//#define _DEBUG
+#define _DEBUG_PREFIX "tt-"
+#include "cuda_debug.cuh"
+
+void quark_blake512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_outputHash, int order);
+
+static bool init[MAX_GPUS] = { 0 };
+
+extern "C" int scanhash_bitcore(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done)
+{
+ uint32_t *pdata = work->data;
+ uint32_t *ptarget = work->target;
+ const uint32_t first_nonce = pdata[19];
+ int intensity = (device_sm[device_map[thr_id]] >= 500 && !is_windows()) ? 20 : 19;
+ uint32_t throughput = cuda_default_throughput(thr_id, 1U << intensity); // 19=256*256*8;
+ //if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);
+
+ if (opt_benchmark) pdata[17] = swab32(0x59090909);
+
+ if (opt_debug || s_ntime != pdata[17] || s_sequence == UINT32_MAX) {
+ uint32_t ntime = swab32(work->data[17]);
+ get_travel_order(ntime, hashOrder);
+ s_ntime = pdata[17];
+ if (opt_debug && !thr_id) {
+ applog(LOG_DEBUG, "timetravel10 hash order %s (%08x)", hashOrder, ntime);
+ }
+ }
+
+ if (opt_benchmark)
+ ptarget[7] = 0x5;
+
+ if (!init[thr_id])
+ {
+ cudaSetDevice(device_map[thr_id]);
+ if (opt_cudaschedule == -1 && gpu_threads == 1) {
+ cudaDeviceReset();
+ // reduce cpu usage
+ cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
+ CUDA_LOG_ERROR();
+ }
+ gpulog(LOG_INFO, thr_id, "Intensity set to %g, %u cuda threads", throughput2intensity(throughput), throughput);
+
+ quark_blake512_cpu_init(thr_id, throughput);
+ quark_bmw512_cpu_init(thr_id, throughput);
+ quark_groestl512_cpu_init(thr_id, throughput);
+ quark_skein512_cpu_init(thr_id, throughput);
+ quark_keccak512_cpu_init(thr_id, throughput);
+ quark_jh512_cpu_init(thr_id, throughput);
+ x11_luffa512_cpu_init(thr_id, throughput);
+ x11_cubehash512_cpu_init(thr_id, throughput);
+ x11_shavite512_cpu_init(thr_id, throughput);
+ if (x11_simd512_cpu_init(thr_id, throughput) != 0) {
+ return 0;
+ }
+#if HASH_FUNC_COUNT > 10
+ x11_echo512_cpu_init(thr_id, throughput);
+#endif
+ CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], (size_t) 64 * throughput), -1);
+ CUDA_CALL_OR_RET_X(cudaMemset(d_hash[thr_id], 0, (size_t) 64 * throughput), -1);
+
+ cuda_check_cpu_init(thr_id, throughput);
+
+ init[thr_id] = true;
+ }
+
+ uint32_t endiandata[20];
+ for (int k=0; k < 19; k++)
+ be32enc(&endiandata[k], pdata[k]);
+
+ cuda_check_cpu_setTarget(ptarget);
+
+ const int hashes = (int) strlen(hashOrder);
+ const char first = hashOrder[0];
+ const uint8_t algo80 = first >= 'A' ? first - 'A' + 10 : first - '0';
+ if (algo80 != s_firstalgo) {
+ s_firstalgo = algo80;
+ }
+
+ // first algo seems locked to blake in bitcore, fine!
+ quark_blake512_cpu_setBlock_80(thr_id, endiandata);
+
+ do {
+ // Hash with CUDA
+
+ quark_blake512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]);
+ TRACE("blake80:");
+
+ for (int i = 1; i < hashes; i++)
+ {
+ const char elem = hashOrder[i];
+ const uint8_t algo64 = elem >= 'A' ? elem - 'A' + 10 : elem - '0';
+
+ switch (algo64) {
+ case BLAKE:
+ quark_blake512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], i);
+ TRACE("blake :");
+ break;
+ case BMW:
+ quark_bmw512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], i);
+ TRACE("bmw :");
+ break;
+ case GROESTL:
+ quark_groestl512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], i);
+ TRACE("groestl:");
+ break;
+ case SKEIN:
+ quark_skein512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], i);
+ TRACE("skein :");
+ break;
+ case JH:
+ quark_jh512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], i);
+ TRACE("jh512 :");
+ break;
+ case KECCAK:
+ quark_keccak512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], i);
+ TRACE("keccak :");
+ break;
+ case LUFFA:
+ x11_luffa512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], i);
+ TRACE("luffa :");
+ break;
+ case CUBEHASH:
+ x11_cubehash512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], i);
+ TRACE("cube :");
+ break;
+ case SHAVITE:
+ x11_shavite512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], i);
+ TRACE("shavite:");
+ break;
+ case SIMD:
+ x11_simd512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], i);
+ TRACE("simd :");
+ break;
+#if HASH_FUNC_COUNT > 10
+ case ECHO:
+ x11_echo512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], i);
+ TRACE("echo :");
+ break;
+#endif
+ }
+ }
+
+ *hashes_done = pdata[19] - first_nonce + throughput;
+
+ work->nonces[0] = cuda_check_hash(thr_id, throughput, pdata[19], d_hash[thr_id]);
+ if (work->nonces[0] != UINT32_MAX)
+ {
+ uint32_t _ALIGN(64) vhash[8];
+ const uint32_t Htarg = ptarget[7];
+ be32enc(&endiandata[19], work->nonces[0]);
+ bitcore_hash(vhash, endiandata);
+
+ if (vhash[7] <= Htarg && fulltest(vhash, ptarget)) {
+ work->valid_nonces = 1;
+ work_set_target_ratio(work, vhash);
+ work->nonces[1] = cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], 1);
+ pdata[19] = work->nonces[0];
+ if (work->nonces[1] != 0) {
+ be32enc(&endiandata[19], work->nonces[1]);
+ bitcore_hash(vhash, endiandata);
+ if (vhash[7] <= Htarg && fulltest(vhash, ptarget)) {
+ bn_set_target_ratio(work, vhash, 1);
+ work->valid_nonces++;
+ }
+ pdata[19] = max(pdata[19], work->nonces[1]) + 1;
+ }
+ return work->valid_nonces;
+ } else if (vhash[7] > Htarg) {
+ gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", work->nonces[0]);
+ pdata[19] = work->nonces[0] + 1;
+ continue;
+ }
+ }
+
+ if ((uint64_t) throughput + pdata[19] >= max_nonce) {
+ pdata[19] = max_nonce;
+ break;
+ }
+ pdata[19] += throughput;
+
+ } while (!work_restart[thr_id].restart);
+
+ *hashes_done = pdata[19] - first_nonce;
+ return 0;
+}
+
+// cleanup
+extern "C" void free_bitcore(int thr_id)
+{
+ if (!init[thr_id])
+ return;
+
+ cudaThreadSynchronize();
+
+ cudaFree(d_hash[thr_id]);
+
+ quark_blake512_cpu_free(thr_id);
+ quark_groestl512_cpu_free(thr_id);
+ x11_simd512_cpu_free(thr_id);
+
+ cuda_check_cpu_free(thr_id);
+ init[thr_id] = false;
+
+ cudaDeviceSynchronize();
+}