forked from tpruvot/ccminer
-
Notifications
You must be signed in to change notification settings - Fork 132
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
- Loading branch information
1 parent
fd406f0
commit e216405
Showing
7 changed files
with
266 additions
and
7 deletions.
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,242 @@ | ||
| /* | ||
| * X13 algorithm | ||
| */ | ||
| 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" | ||
| #include "sph/sph_echo.h" | ||
|
|
||
| #include "sph/sph_hamsi.h" | ||
| #include "sph/sph_fugue.h" | ||
| } | ||
| #include "miner.h" | ||
|
|
||
| #include "cuda_helper.h" | ||
| #include "x11/cuda_x11.h" | ||
|
|
||
| static uint32_t *d_hash[MAX_GPUS]; | ||
|
|
||
| extern void x13_hamsi512_cpu_init(int thr_id, uint32_t threads); | ||
| extern void x13_hamsi512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); | ||
|
|
||
| // X13 CPU Hash | ||
| extern "C" void x12hash(void *output, const void *input) | ||
| { | ||
| // blake1-bmw2-grs3-skein4-jh5-keccak6-luffa7-cubehash8-shavite9-simd10-echo11-hamsi12-fugue13 | ||
|
|
||
| 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; | ||
|
|
||
| uint32_t hash[32]; | ||
| memset(hash, 0, sizeof hash); | ||
|
|
||
| sph_blake512_init(&ctx_blake); | ||
| sph_blake512 (&ctx_blake, input, 80); | ||
| sph_blake512_close(&ctx_blake, (void*) hash); | ||
|
|
||
| sph_bmw512_init(&ctx_bmw); | ||
| sph_bmw512 (&ctx_bmw, (const void*) hash, 64); | ||
| sph_bmw512_close(&ctx_bmw, (void*) hash); | ||
|
|
||
| sph_luffa512_init(&ctx_luffa); | ||
| sph_luffa512(&ctx_luffa, (const void*)hash, 64); | ||
| sph_luffa512_close(&ctx_luffa, (void*)hash); | ||
|
|
||
| sph_cubehash512_init(&ctx_cubehash); | ||
| sph_cubehash512(&ctx_cubehash, (const void*)hash, 64); | ||
| sph_cubehash512_close(&ctx_cubehash, (void*)hash); | ||
|
|
||
| sph_shavite512_init(&ctx_shavite); | ||
| sph_shavite512(&ctx_shavite, (const void*)hash, 64); | ||
| sph_shavite512_close(&ctx_shavite, (void*)hash); | ||
|
|
||
| sph_simd512_init(&ctx_simd); | ||
| sph_simd512(&ctx_simd, (const void*)hash, 64); | ||
| sph_simd512_close(&ctx_simd, (void*)hash); | ||
|
|
||
| sph_echo512_init(&ctx_echo); | ||
| sph_echo512(&ctx_echo, (const void*)hash, 64); | ||
| sph_echo512_close(&ctx_echo, (void*)hash); | ||
|
|
||
| sph_groestl512_init(&ctx_groestl); | ||
| sph_groestl512 (&ctx_groestl, (const void*) hash, 64); | ||
| sph_groestl512_close(&ctx_groestl, (void*) hash); | ||
|
|
||
| sph_skein512_init(&ctx_skein); | ||
| sph_skein512 (&ctx_skein, (const void*) hash, 64); | ||
| sph_skein512_close(&ctx_skein, (void*) hash); | ||
|
|
||
| sph_jh512_init(&ctx_jh); | ||
| sph_jh512 (&ctx_jh, (const void*) hash, 64); | ||
| 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_hamsi512_init(&ctx_hamsi); | ||
| sph_hamsi512 (&ctx_hamsi, (const void*) hash, 64); | ||
| sph_hamsi512_close(&ctx_hamsi, (void*) hash); | ||
|
|
||
| memcpy(output, hash, 32); | ||
| } | ||
|
|
||
| static bool init[MAX_GPUS] = { 0 }; | ||
|
|
||
| extern "C" int scanhash_x12(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 = 19; // (device_sm[device_map[thr_id]] > 500 && !is_windows()) ? 20 : 19; | ||
| uint32_t throughput = cuda_default_throughput(thr_id, 1 << intensity); // 19=256*256*8; | ||
| //if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce); | ||
|
|
||
| if (opt_benchmark) | ||
| ((uint32_t*)ptarget)[7] = 0x000f; | ||
|
|
||
| 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); | ||
| x11_luffaCubehash512_cpu_init(thr_id, throughput); | ||
| x11_shavite512_cpu_init(thr_id, throughput); | ||
| if (x11_simd512_cpu_init(thr_id, throughput) != 0) { | ||
| return 0; | ||
| } | ||
| x11_echo512_cpu_init(thr_id, throughput); | ||
| quark_groestl512_cpu_init(thr_id, throughput); | ||
| quark_skein512_cpu_init(thr_id, throughput); | ||
| quark_bmw512_cpu_init(thr_id, throughput); | ||
| quark_keccak512_cpu_init(thr_id, throughput); | ||
| quark_jh512_cpu_init(thr_id, throughput); | ||
| x13_hamsi512_cpu_init(thr_id, throughput); | ||
|
|
||
| CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput), 0); | ||
|
|
||
| cuda_check_cpu_init(thr_id, throughput); | ||
|
|
||
| init[thr_id] = true; | ||
| } | ||
|
|
||
| uint32_t endiandata[20]; | ||
| for (int k=0; k < 20; k++) | ||
| be32enc(&endiandata[k], pdata[k]); | ||
|
|
||
| quark_blake512_cpu_setBlock_80(thr_id, endiandata); | ||
| cuda_check_cpu_setTarget(ptarget); | ||
|
|
||
| do { | ||
| int order = 0; | ||
|
|
||
| quark_blake512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]); order++; | ||
| quark_bmw512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
| x11_luffaCubehash512_cpu_hash_64(thr_id, throughput, d_hash[thr_id], order++); | ||
| x11_shavite512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
| x11_simd512_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++); | ||
| quark_groestl512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
| quark_skein512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
| quark_jh512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
| quark_keccak512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
| x13_hamsi512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
|
|
||
| *hashes_done = pdata[19] - first_nonce + throughput; | ||
|
|
||
| CUDA_LOG_ERROR(); | ||
|
|
||
| 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]); | ||
| x12hash(vhash, endiandata); | ||
|
|
||
| if (vhash[7] <= ptarget[7] && 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]); | ||
| x13hash(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 | ||
| } | ||
| return work->valid_nonces; | ||
| } | ||
| 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); | ||
|
|
||
| *hashes_done = pdata[19] - first_nonce; | ||
|
|
||
| CUDA_LOG_ERROR(); | ||
|
|
||
| return 0; | ||
| } | ||
|
|
||
| // cleanup | ||
| extern "C" void free_x12(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); | ||
| CUDA_LOG_ERROR(); | ||
|
|
||
| cudaDeviceSynchronize(); | ||
| init[thr_id] = false; | ||
| } |