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.
only made a few changes to keep algos in the alpha order
- Loading branch information
Showing
12 changed files
with
227 additions
and
6 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
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,197 @@ | ||
| // | ||
| // | ||
| // PHI1612 algo | ||
| // Skein + JH + CubeHash + Fugue + Gost + Echo | ||
| // | ||
| // Implemented by anorganix @ bitcointalk on 01.10.2017 | ||
| // Feel free to send some satoshis to 1Bitcoin8tfbtGAQNFxDRUVUfFgFWKoWi9 | ||
| // | ||
| // | ||
|
|
||
| extern "C" { | ||
| #include "sph/sph_skein.h" | ||
| #include "sph/sph_jh.h" | ||
| #include "sph/sph_cubehash.h" | ||
| #include "sph/sph_fugue.h" | ||
| #include "sph/sph_streebog.h" | ||
| #include "sph/sph_echo.h" | ||
| } | ||
|
|
||
| #include "miner.h" | ||
| #include "cuda_helper.h" | ||
| #include "cuda_x11.h" | ||
|
|
||
| extern void skein512_cpu_setBlock_80(void *pdata); | ||
| extern void skein512_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_hash, int swap); | ||
| extern void streebog_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash); | ||
|
|
||
| extern void x13_fugue512_cpu_init(int thr_id, uint32_t threads); | ||
| extern void x13_fugue512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); | ||
| extern void x13_fugue512_cpu_free(int thr_id); | ||
|
|
||
| #include <stdio.h> | ||
| #include <memory.h> | ||
|
|
||
| static uint32_t *d_hash[MAX_GPUS]; | ||
|
|
||
| extern "C" void phihash(void *output, const void *input) | ||
| { | ||
| unsigned char _ALIGN(128) hash[128] = { 0 }; | ||
|
|
||
| sph_skein512_context ctx_skein; | ||
| sph_jh512_context ctx_jh; | ||
| sph_cubehash512_context ctx_cubehash; | ||
| sph_fugue512_context ctx_fugue; | ||
| sph_gost512_context ctx_gost; | ||
| sph_echo512_context ctx_echo; | ||
|
|
||
| sph_skein512_init(&ctx_skein); | ||
| sph_skein512(&ctx_skein, input, 80); | ||
| 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_cubehash512_init(&ctx_cubehash); | ||
| sph_cubehash512(&ctx_cubehash, (const void*)hash, 64); | ||
| sph_cubehash512_close(&ctx_cubehash, (void*)hash); | ||
|
|
||
| sph_fugue512_init(&ctx_fugue); | ||
| sph_fugue512(&ctx_fugue, (const void*)hash, 64); | ||
| sph_fugue512_close(&ctx_fugue, (void*)hash); | ||
|
|
||
| sph_gost512_init(&ctx_gost); | ||
| sph_gost512(&ctx_gost, (const void*)hash, 64); | ||
| sph_gost512_close(&ctx_gost, (void*)hash); | ||
|
|
||
| sph_echo512_init(&ctx_echo); | ||
| sph_echo512(&ctx_echo, (const void*)hash, 64); | ||
| sph_echo512_close(&ctx_echo, (void*)hash); | ||
|
|
||
| memcpy(output, hash, 32); | ||
| } | ||
|
|
||
| #define _DEBUG_PREFIX "phi" | ||
| #include "cuda_debug.cuh" | ||
|
|
||
| static bool init[MAX_GPUS] = { 0 }; | ||
|
|
||
| extern "C" int scanhash_phi(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()) ? 19 : 18; // 2^18 = 262144 cuda threads | ||
| if (device_sm[dev_id] >= 600) 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) | ||
| ptarget[7] = 0xf; | ||
|
|
||
| if (!init[thr_id]) | ||
| { | ||
| cudaSetDevice(device_map[thr_id]); | ||
| if (opt_cudaschedule == -1 && gpu_threads == 1) { | ||
| cudaDeviceReset(); | ||
| cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); | ||
|
|
||
| CUDA_LOG_ERROR(); | ||
| } | ||
| gpulog(LOG_INFO, thr_id, "Intensity set to %g, %u cuda threads", throughput2intensity(throughput), throughput); | ||
|
|
||
| quark_skein512_cpu_init(thr_id, throughput); | ||
| quark_jh512_cpu_init(thr_id, throughput); | ||
| x11_cubehash512_cpu_init(thr_id, throughput); | ||
| x13_fugue512_cpu_init(thr_id, throughput); | ||
| x11_echo512_cpu_init(thr_id, throughput); | ||
|
|
||
| CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], (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 < 20; k++) | ||
| be32enc(&endiandata[k], pdata[k]); | ||
|
|
||
| skein512_cpu_setBlock_80((void*)endiandata); | ||
| cuda_check_cpu_setTarget(ptarget); | ||
|
|
||
| do { | ||
| int order = 0; | ||
|
|
||
| skein512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], 1); order++; | ||
| quark_jh512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
| x11_cubehash512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
| x13_fugue512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
| streebog_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); | ||
| x11_echo512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
|
|
||
| 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]); | ||
| phihash(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); | ||
| *hashes_done = pdata[19] - first_nonce + throughput; | ||
| if (work->nonces[1] != 0) { | ||
| be32enc(&endiandata[19], work->nonces[1]); | ||
| phihash(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; | ||
| return 0; | ||
| } | ||
|
|
||
| // cleanup | ||
| extern "C" void free_phi(int thr_id) | ||
| { | ||
| if (!init[thr_id]) | ||
| return; | ||
|
|
||
| cudaThreadSynchronize(); | ||
| cudaFree(d_hash[thr_id]); | ||
|
|
||
| cuda_check_cpu_free(thr_id); | ||
| init[thr_id] = false; | ||
|
|
||
| cudaDeviceSynchronize(); | ||
| } |