forked from tpruvot/ccminer
-
Notifications
You must be signed in to change notification settings - Fork 1
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Signed-off-by: Tanguy Pruvot <[email protected]>
- Loading branch information
Showing
10 changed files
with
188 additions
and
1 deletion.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,166 @@ | ||
| /** | ||
| * Tribus Algo for Denarius | ||
| * | ||
| * tpruvot@github 06 2017 - GPLv3 | ||
| * | ||
| */ | ||
| extern "C" { | ||
| #include "sph/sph_jh.h" | ||
| #include "sph/sph_keccak.h" | ||
| #include "sph/sph_echo.h" | ||
| } | ||
|
|
||
| #include "miner.h" | ||
| #include "cuda_helper.h" | ||
| #include "x11/cuda_x11.h" | ||
|
|
||
| void jh512_setBlock_80(int thr_id, uint32_t *endiandata); | ||
| void jh512_cuda_hash_80(const int thr_id, const uint32_t threads, const uint32_t startNounce, uint32_t *d_hash); | ||
|
|
||
| static uint32_t *d_hash[MAX_GPUS]; | ||
|
|
||
|
|
||
| // cpu hash | ||
|
|
||
| extern "C" void tribus_hash(void *state, const void *input) | ||
| { | ||
| uint8_t _ALIGN(64) hash[64]; | ||
|
|
||
| sph_jh512_context ctx_jh; | ||
| sph_keccak512_context ctx_keccak; | ||
| sph_echo512_context ctx_echo; | ||
|
|
||
| sph_jh512_init(&ctx_jh); | ||
| sph_jh512(&ctx_jh, input, 80); | ||
| sph_jh512_close(&ctx_jh, (void*) hash); | ||
|
|
||
| sph_keccak512_init(&ctx_keccak); | ||
| sph_keccak512(&ctx_keccak, (const void*) hash, 64); | ||
| sph_keccak512_close(&ctx_keccak, (void*) hash); | ||
|
|
||
| sph_echo512_init(&ctx_echo); | ||
| sph_echo512(&ctx_echo, (const void*) hash, 64); | ||
| sph_echo512_close(&ctx_echo, (void*) hash); | ||
|
|
||
| memcpy(state, hash, 32); | ||
| } | ||
|
|
||
| static bool init[MAX_GPUS] = { 0 }; | ||
|
|
||
| extern "C" int scanhash_tribus(int thr_id, struct work *work, uint32_t max_nonce, unsigned long *hashes_done) | ||
| { | ||
| uint32_t _ALIGN(64) endiandata[20]; | ||
| uint32_t *pdata = work->data; | ||
| uint32_t *ptarget = work->target; | ||
| const uint32_t first_nonce = pdata[19]; | ||
|
|
||
| int8_t intensity = is_windows() ? 20 : 23; | ||
| uint32_t throughput = cuda_default_throughput(thr_id, 1 << intensity); | ||
| if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce); | ||
|
|
||
| if (opt_benchmark) | ||
| ((uint32_t*)ptarget)[7] = 0x00FF; | ||
|
|
||
| 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_jh512_cpu_init(thr_id, throughput); | ||
| quark_keccak512_cpu_init(thr_id, throughput); | ||
| x11_echo512_cpu_init(thr_id, throughput); | ||
|
|
||
| // char[64] work space for hashes results | ||
| CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], (size_t)64 * throughput)); | ||
|
|
||
| cuda_check_cpu_init(thr_id, throughput); | ||
| init[thr_id] = true; | ||
| } | ||
|
|
||
| for (int k=0; k < 20; k++) | ||
| be32enc(&endiandata[k], pdata[k]); | ||
|
|
||
| jh512_setBlock_80(thr_id, endiandata); | ||
| cuda_check_cpu_setTarget(ptarget); | ||
|
|
||
| work->valid_nonces = 0; | ||
|
|
||
| do { | ||
| int order = 1; | ||
|
|
||
| // Hash with CUDA | ||
| jh512_cuda_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]); | ||
| quark_keccak512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
| x11_echo512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
|
|
||
| *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) | ||
| { | ||
| const uint32_t Htarg = ptarget[7]; | ||
| uint32_t _ALIGN(64) vhash[8]; | ||
| be32enc(&endiandata[19], work->nonces[0]); | ||
| tribus_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); | ||
| if (work->nonces[1] != 0) { | ||
| be32enc(&endiandata[19], work->nonces[1]); | ||
| tribus_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 | ||
| } | ||
| goto out; | ||
| } | ||
| else if (vhash[7] > Htarg) { | ||
| gpu_increment_reject(thr_id); | ||
| if (!opt_quiet) | ||
| 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); | ||
|
|
||
| out: | ||
| // *hashes_done = pdata[19] - first_nonce; | ||
|
|
||
| return work->valid_nonces; | ||
| } | ||
|
|
||
| // ressources cleanup | ||
| extern "C" void free_tribus(int thr_id) | ||
| { | ||
| if (!init[thr_id]) | ||
| return; | ||
|
|
||
| cudaThreadSynchronize(); | ||
|
|
||
| cudaFree(d_hash[thr_id]); | ||
|
|
||
| quark_groestl512_cpu_free(thr_id); | ||
| cuda_check_cpu_free(thr_id); | ||
| init[thr_id] = false; | ||
|
|
||
| cudaDeviceSynchronize(); | ||
| } |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters