diff --git a/algos.h b/algos.h index 2cd1f2bfd9..ea8a469ac6 100644 --- a/algos.h +++ b/algos.h @@ -66,6 +66,7 @@ enum sha_algos { ALGO_X14, ALGO_X15, ALGO_X16R, + ALGO_X16RT, ALGO_X16RV2, ALGO_X16S, ALGO_X17, @@ -145,6 +146,7 @@ static const char *algo_names[] = { "x14", "x15", "x16r", + "x16rt", "x16rv2", "x16s", "x17", diff --git a/ccminer.cpp b/ccminer.cpp index 87b7a3d1e0..d5028a39f6 100644 --- a/ccminer.cpp +++ b/ccminer.cpp @@ -301,6 +301,7 @@ Options:\n\ x15 X15\n\ x17 X17\n\ x16r X16R\n\ + x16rt X16RT\n\ x16rv2 X16R V2\n\ x16s X16S\n\ x21s X21S\n\ @@ -1747,6 +1748,7 @@ static bool stratum_gen_work(struct stratum_ctx *sctx, struct work *work) case ALGO_BITCORE: case ALGO_BMW512: case ALGO_X16R: + case ALGO_X16RT: case ALGO_X16RV2: case ALGO_X16S: case ALGO_X21S: @@ -2559,6 +2561,9 @@ static void *miner_thread(void *userdata) case ALGO_X16R: rc = scanhash_x16r(thr_id, &work, max_nonce, &hashes_done); break; + case ALGO_X16RT: + rc = scanhash_x16rt(thr_id, &work, max_nonce, &hashes_done); + break; case ALGO_X16RV2: rc = scanhash_x16rv2(thr_id, &work, max_nonce, &hashes_done); break; diff --git a/miner.h b/miner.h index dbab589f22..f307687c72 100644 --- a/miner.h +++ b/miner.h @@ -330,6 +330,7 @@ extern int scanhash_x13(int thr_id, struct work* work, uint32_t max_nonce, unsig extern int scanhash_x14(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done); extern int scanhash_x15(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done); extern int scanhash_x16r(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done); +extern int scanhash_x16rt(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done); extern int scanhash_x16rv2(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done); extern int scanhash_x16s(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done); extern int scanhash_x17(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done); @@ -396,6 +397,7 @@ extern void free_x13(int thr_id); extern void free_x14(int thr_id); extern void free_x15(int thr_id); extern void free_x16r(int thr_id); +extern void free_x16rt(int thr_id); extern void free_x16rv2(int thr_id); extern void free_x16s(int thr_id); extern void free_x17(int thr_id); @@ -945,6 +947,7 @@ void x13hash(void *output, const void *input); void x14hash(void *output, const void *input); void x15hash(void *output, const void *input); void x16r_hash(void *output, const void *input); +void x16rt_hash(void *output, const void *input); void x16rv2_hash(void *output, const void *input); void x16s_hash(void *output, const void *input); void x17hash(void *output, const void *input); diff --git a/util.cpp b/util.cpp index 03dec94dd8..00e5dc54f8 100644 --- a/util.cpp +++ b/util.cpp @@ -2333,6 +2333,9 @@ void print_hash_tests(void) x16r_hash(&hash[0], &buf[0]); printpfx("X16R", hash); + x16rt_hash(&hash[0], &buf[0]); + printpfx("X16RT", hash); + x16rv2_hash(&hash[0], &buf[0]); printpfx("X16Rv2", hash); diff --git a/x16/x16r.cu b/x16/x16r.cu index 2caa5d0f1b..24ed362d5a 100644 --- a/x16/x16r.cu +++ b/x16/x16r.cu @@ -621,3 +621,524 @@ extern "C" void free_x16r(int thr_id) cudaDeviceSynchronize(); init[thr_id] = false; } + +static void x16rt_getTimeHash( const uint32_t timeStamp, void* timeHash ) +{ + int32_t maskedTime = timeStamp & 0xffffff80; + sha256d( (unsigned char*)timeHash, (const unsigned char*)( &maskedTime ), + sizeof( maskedTime ) ); +} + +static void x16rt_getAlgoString( const uint32_t *timeHash, char *output) +{ + char *sptr = output; + uint8_t* data = (uint8_t*)timeHash; + + for (uint8_t j = 0; j < HASH_FUNC_COUNT; j++) { + uint8_t b = (15 - j) >> 1; // 16 ascii hex chars, reversed + uint8_t algoDigit = (j & 1) ? data[b] & 0xF : data[b] >> 4; + + if (algoDigit >= 10) + sprintf(sptr, "%c", 'A' + (algoDigit - 10)); + else + sprintf(sptr, "%u", (uint32_t) algoDigit); + sptr++; + } + *sptr = '\0'; +} + +extern "C" void x16rt_hash(void *output, const void *input) +{ + unsigned char _ALIGN(64) hash[128]; + uint32_t _ALIGN(64) timeHash[8]; + + sph_blake512_context ctx_blake; + sph_bmw512_context ctx_bmw; + sph_groestl512_context ctx_groestl; + sph_jh512_context ctx_jh; + sph_keccak512_context ctx_keccak; + sph_skein512_context ctx_skein; + sph_luffa512_context ctx_luffa; + sph_cubehash512_context ctx_cubehash; + sph_shavite512_context ctx_shavite; + sph_simd512_context ctx_simd; + sph_echo512_context ctx_echo; + sph_hamsi512_context ctx_hamsi; + sph_fugue512_context ctx_fugue; + sph_shabal512_context ctx_shabal; + sph_whirlpool_context ctx_whirlpool; + sph_sha512_context ctx_sha512; + + void *in = (void*) input; + int size = 80; + + uint32_t *in32 = (uint32_t*) input; + x16rt_getTimeHash( in32[17], &timeHash ); + x16rt_getAlgoString( &timeHash[0], hashOrder ); + char elem = hashOrder[0]; + uint8_t algo = elem >= 'A' ? elem - 'A' + 10 : elem - '0'; + + for (int i = 0; i < 16; i++) + { + 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_luffa); + sph_luffa512(&ctx_luffa, in, size); + sph_luffa512_close(&ctx_luffa, hash); + break; + case CUBEHASH: + sph_cubehash512_init(&ctx_cubehash); + sph_cubehash512(&ctx_cubehash, in, size); + sph_cubehash512_close(&ctx_cubehash, hash); + break; + case SHAVITE: + sph_shavite512_init(&ctx_shavite); + sph_shavite512(&ctx_shavite, in, size); + sph_shavite512_close(&ctx_shavite, hash); + break; + case SIMD: + sph_simd512_init(&ctx_simd); + sph_simd512(&ctx_simd, in, size); + sph_simd512_close(&ctx_simd, hash); + break; + case ECHO: + sph_echo512_init(&ctx_echo); + sph_echo512(&ctx_echo, in, size); + sph_echo512_close(&ctx_echo, hash); + break; + case HAMSI: + sph_hamsi512_init(&ctx_hamsi); + sph_hamsi512(&ctx_hamsi, in, size); + sph_hamsi512_close(&ctx_hamsi, hash); + break; + case FUGUE: + sph_fugue512_init(&ctx_fugue); + sph_fugue512(&ctx_fugue, in, size); + sph_fugue512_close(&ctx_fugue, hash); + break; + case SHABAL: + sph_shabal512_init(&ctx_shabal); + sph_shabal512(&ctx_shabal, in, size); + sph_shabal512_close(&ctx_shabal, hash); + break; + case WHIRLPOOL: + sph_whirlpool_init(&ctx_whirlpool); + sph_whirlpool(&ctx_whirlpool, in, size); + sph_whirlpool_close(&ctx_whirlpool, hash); + break; + case SHA512: + sph_sha512_init(&ctx_sha512); + sph_sha512(&ctx_sha512,(const void*) in, size); + sph_sha512_close(&ctx_sha512,(void*) hash); + break; + } + in = (void*) hash; + size = 64; + + if (i!=15) + { + elem = hashOrder[i+1]; + algo = elem >= 'A' ? elem - 'A' + 10 : elem - '0'; + } + + } + memcpy(output, hash, 32); +} + +extern "C" int scanhash_x16rt(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]; + const int dev_id = device_map[thr_id]; + int intensity = (device_sm[dev_id] > 500 && !is_windows()) ? 20 : 19; + if (strstr(device_name[dev_id], "GTX 1080")) intensity = 20; + uint32_t throughput = cuda_default_throughput(thr_id, 1U << intensity); + //if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce); + if (opt_benchmark) + { + ((uint32_t*)ptarget)[7] = 0x0ff; + ((uint32_t*)pdata)[1] = 0x88888888; + ((uint32_t*)pdata)[2] = 0x88888888; + } + + uint32_t _ALIGN(64) endiandata[20]; + uint32_t _ALIGN(64) timeHash[8]; + + for (int k = 0; k < 19; k++) + be32enc(&endiandata[k], pdata[k]); + + uint32_t masked_ntime = endiandata[17] & 0xffffff80; + if ( s_ntime != masked_ntime ) + { + x16rt_getTimeHash( masked_ntime, &timeHash ); + x16rt_getAlgoString( &timeHash[0], hashOrder ); + s_ntime = masked_ntime; + if ( !(thr_id || opt_quiet) ) + { + applog( LOG_INFO, "hash order %s (%08x)", hashOrder, + endiandata[17] ); + cudaDeviceSynchronize(); + } + } + + + if (!init[thr_id]) + { + cudaSetDevice(dev_id); + if (opt_cudaschedule == -1 && gpu_threads == 1) { + cudaDeviceReset(); + // reduce cpu usage + cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); + } + gpulog(LOG_INFO, thr_id, "Intensity set to %g, %u cuda threads", throughput2intensity(throughput), throughput); + + cuda_get_arch(thr_id); + use_compat_kernels[thr_id] = (cuda_arch[dev_id] < 500); + if (use_compat_kernels[thr_id]) + x11_echo512_cpu_init(thr_id, 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_jh512_cpu_init(thr_id, throughput); + quark_keccak512_cpu_init(thr_id, throughput); + x11_shavite512_cpu_init(thr_id, throughput); + x11_simd512_cpu_init(thr_id, (throughput)); + x13_hamsi512_cpu_init(thr_id, throughput); + x16_fugue512_cpu_init(thr_id, throughput); + x15_whirlpool_cpu_init(thr_id, throughput, 0); + x16_whirlpool512_init(thr_id, throughput); + x11_luffa512_cpu_init(thr_id, throughput); // 64 + x16_echo512_cuda_init(thr_id, throughput); + x13_fugue512_cpu_init(thr_id, throughput); + x16_fugue512_cpu_init(thr_id, throughput); + x14_shabal512_cpu_init(thr_id, throughput); + + CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], (size_t) 64 * throughput), 0); + + cuda_check_cpu_init(thr_id, throughput); + + init[thr_id] = true; + } + + cuda_check_cpu_setTarget(ptarget); + + char elem = hashOrder[0]; + const uint8_t algo80 = elem >= 'A' ? elem - 'A' + 10 : elem - '0'; + + switch (algo80) + { + case BLAKE: + quark_blake512_cpu_setBlock_80(thr_id, endiandata); + break; + case BMW: + quark_bmw512_cpu_setBlock_80(endiandata); + break; + case GROESTL: + groestl512_setBlock_80(thr_id, endiandata); + break; + case JH: + jh512_setBlock_80(thr_id, endiandata); + break; + case KECCAK: + keccak512_setBlock_80(thr_id, endiandata); + break; + case SKEIN: + skein512_cpu_setBlock_80((void*)endiandata); + break; + case LUFFA: + qubit_luffa512_cpu_setBlock_80((void*)endiandata); + break; + case CUBEHASH: + cubehash512_setBlock_80(thr_id, endiandata); + break; + case SHAVITE: + x11_shavite512_setBlock_80((void*)endiandata); + break; + case SIMD: + x16_simd512_setBlock_80((void*)endiandata); + break; + case ECHO: + x16_echo512_setBlock_80((void*)endiandata); + break; + case HAMSI: + x16_hamsi512_setBlock_80((uint64_t*)endiandata); + break; + case FUGUE: + x16_fugue512_setBlock_80((void*)pdata); + break; + case SHABAL: + x16_shabal512_setBlock_80((void*)endiandata); + break; + case WHIRLPOOL: + x16_whirlpool512_setBlock_80((void*)endiandata); + break; + case SHA512: + x16_sha512_setBlock_80(endiandata); + break; + default: { + if (!thr_id) + applog(LOG_WARNING, "kernel %s %c unimplemented, order %s", algo_strings[algo80], elem, hashOrder); + sleep(5); + return -1; + } + } + + + int warn = 0; + + do { + int order = 0; + + uint32_t start = pdata[19]; + //uint32_t foundNonce; + + // Hash with CUDA + switch (algo80) + { + case BLAKE: + quark_blake512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]); order++; + break; + case BMW: + quark_bmw512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++); + break; + case GROESTL: + groestl512_cuda_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]); order++; + break; + case JH: + jh512_cuda_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]); order++; + break; + case KECCAK: + keccak512_cuda_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]); order++; + break; + case SKEIN: + skein512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], 1); order++; + break; + case LUFFA: + qubit_luffa512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++); + break; + case CUBEHASH: + cubehash512_cuda_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]); order++; + break; + case SHAVITE: + x11_shavite512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++); + break; + case SIMD: + x16_simd512_cuda_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]); order++; + break; + case ECHO: + x16_echo512_cuda_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]); order++; + break; + case HAMSI: + x16_hamsi512_cuda_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]); order++; + break; + case FUGUE: + x16_fugue512_cuda_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]); order++; + break; + case SHABAL: + x16_shabal512_cuda_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]); order++; + break; + case WHIRLPOOL: + x16_whirlpool512_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]); order++; + break; + case SHA512: + x16_sha512_cuda_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]); order++; + break; + } + uint8_t algo; + + + for (int i = 1; i < 16; i++) + { + elem = hashOrder[i]; + 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], order++); + break; + case BMW: + quark_bmw512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + break; + case GROESTL: + quark_groestl512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + break; + case JH: + quark_jh512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + break; + case KECCAK: + quark_keccak512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + break; + case SKEIN: + quark_skein512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + break; + case LUFFA: + x11_luffa512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + break; + case CUBEHASH: + x11_cubehash512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + break; + case SHAVITE: + x11_shavite512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + break; + case SIMD: + x11_simd512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + break; + case ECHO: + if (use_compat_kernels[thr_id]) + x11_echo512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + else { + x16_echo512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); order++; + } + break; + case HAMSI: + x13_hamsi512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + break; + case FUGUE: + x13_fugue512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + break; + case SHABAL: + x14_shabal512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + break; + case WHIRLPOOL: + x15_whirlpool_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); + break; + case SHA512: + x17_sha512_cpu_hash_64(thr_id, throughput, pdata[19], d_hash[thr_id]); order++; + break; + } + } + + *hashes_done = pdata[19] - first_nonce + throughput; + work->nonces[0] = cuda_check_hash(thr_id, throughput, pdata[19], d_hash[thr_id]); +#ifdef _DEBUG + uint32_t _ALIGN(64) dhash[8]; + be32enc(&endiandata[19], pdata[19]); + x16rt_hash(dhash, endiandata); + applog_hash(dhash); + return -1; +#endif + if (work->nonces[0] != UINT32_MAX) + { + const uint32_t Htarg = ptarget[7]; + uint32_t _ALIGN(64) vhash[8]; + be32enc(&endiandata[19], work->nonces[0]); + x16rt_hash(vhash, endiandata); + + if (vhash[7] <= Htarg && fulltest(vhash, ptarget)) { + work->valid_nonces = 1; + work->nonces[1] = cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], 1); + work_set_target_ratio(work, vhash); + if (work->nonces[1] != 0) { + be32enc(&endiandata[19], work->nonces[1]); + x16rt_hash(vhash, endiandata); + bn_set_target_ratio(work, vhash, 1); + work->valid_nonces++; + pdata[19] = max(work->nonces[0], work->nonces[1]) + 1; + } else { + pdata[19] = work->nonces[0] + 1; // cursor + } +#if 0 + gpulog(LOG_INFO, thr_id, "hash found with %s 80!", algo_strings[algo80]); + + algo80_tests[algo80] += work->valid_nonces; + char oks64[128] = { 0 }; + char oks80[128] = { 0 }; + char fails[128] = { 0 }; + for (int a = 0; a < HASH_FUNC_COUNT; a++) { + const char elem = hashOrder[a]; + const uint8_t algo64 = elem >= 'A' ? elem - 'A' + 10 : elem - '0'; + if (a > 0) algo64_tests[algo64] += work->valid_nonces; + sprintf(&oks64[strlen(oks64)], "|%X:%2d", a, algo64_tests[a] < 100 ? algo64_tests[a] : 99); + sprintf(&oks80[strlen(oks80)], "|%X:%2d", a, algo80_tests[a] < 100 ? algo80_tests[a] : 99); + sprintf(&fails[strlen(fails)], "|%X:%2d", a, algo80_fails[a] < 100 ? algo80_fails[a] : 99); + } + applog(LOG_INFO, "K64: %s", oks64); + applog(LOG_INFO, "K80: %s", oks80); + applog(LOG_ERR, "F80: %s", fails); +#endif + return work->valid_nonces; + } + else if (vhash[7] > Htarg) { + // x11+ coins could do some random error, but not on retry + gpu_increment_reject(thr_id); + algo80_fails[algo80]++; + if (!warn) { + warn++; + pdata[19] = work->nonces[0] + 1; + continue; + } else { + if (!opt_quiet) gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU! %s %s", + work->nonces[0], algo_strings[algo80], hashOrder); + warn = 0; + } + } + } + + if ((uint64_t)throughput + pdata[19] >= max_nonce) { + pdata[19] = max_nonce; + break; + } + + pdata[19] += throughput; + + } while (pdata[19] < max_nonce && !work_restart[thr_id].restart); + + *hashes_done = pdata[19] - first_nonce; + return 0; +} + +// cleanup +extern "C" void free_x16rt(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); + x13_fugue512_cpu_free(thr_id); + x16_fugue512_cpu_free(thr_id); // to merge with x13_fugue512 ? + x15_whirlpool_cpu_free(thr_id); + + cuda_check_cpu_free(thr_id); + + cudaDeviceSynchronize(); + init[thr_id] = false; +} \ No newline at end of file