forked from cbuchner1/ccminer
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
VeltorCoin Streebog based algo (veltor)
also known as "Thor's Riddle"... yes sure ;) Credits to ocminer who found and "implemented" it. Note: tested "ok" on x64 and CUDA 6.5 x86, not on 7.5 and 8.0 x86 PS: Don't have the time for a more proper CUDA implementation of Streebog
- Loading branch information
Showing
12 changed files
with
209 additions
and
9 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,175 @@ | ||
| extern "C" { | ||
| #include "sph/sph_skein.h" | ||
| #include "sph/sph_shavite.h" | ||
| #include "sph/sph_shabal.h" | ||
| #include "sph/sph_streebog.h" | ||
| } | ||
|
|
||
| #include "miner.h" | ||
| #include "cuda_helper.h" | ||
| #include "cuda_x11.h" | ||
|
|
||
| extern void streebog_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_hash); | ||
|
|
||
| extern void x14_shabal512_cpu_init(int thr_id, uint32_t threads); | ||
| extern void x14_shabal512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); | ||
|
|
||
| extern void quark_skein512_cpu_init(int thr_id, uint32_t threads); | ||
| 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); | ||
|
|
||
| #include <stdio.h> | ||
| #include <memory.h> | ||
|
|
||
| static uint32_t *d_hash[MAX_GPUS]; | ||
|
|
||
| // Veltor CPU Hash | ||
| extern "C" void veltorhash(void *output, const void *input) | ||
| { | ||
| uint8_t _ALIGN(64) hash[128] = { 0 }; | ||
|
|
||
| sph_skein512_context ctx_skein; | ||
| sph_shavite512_context ctx_shavite; | ||
| sph_shabal512_context ctx_shabal; | ||
| sph_gost512_context ctx_gost; | ||
|
|
||
| sph_skein512_init(&ctx_skein); | ||
| sph_skein512(&ctx_skein, input, 80); | ||
| sph_skein512_close(&ctx_skein, (void*) hash); | ||
|
|
||
| sph_shavite512_init(&ctx_shavite); | ||
| sph_shavite512(&ctx_shavite, (const void*) hash, 64); | ||
| sph_shavite512_close(&ctx_shavite, (void*) hash); | ||
|
|
||
| sph_shabal512_init(&ctx_shabal); | ||
| sph_shabal512(&ctx_shabal, (const void*) hash, 64); | ||
| sph_shabal512_close(&ctx_shabal, (void*) hash); | ||
|
|
||
| sph_gost512_init(&ctx_gost); | ||
| sph_gost512(&ctx_gost, (const void*) hash, 64); | ||
| sph_gost512_close(&ctx_gost, (void*) hash); | ||
|
|
||
| memcpy(output, hash, 32); | ||
| } | ||
|
|
||
| //#define _DEBUG | ||
| #define _DEBUG_PREFIX "veltor" | ||
| #include "cuda_debug.cuh" | ||
|
|
||
| static bool init[MAX_GPUS] = { 0 }; | ||
|
|
||
| extern "C" int scanhash_veltor(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 : 18; | ||
| 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) | ||
| ptarget[7] = 0xf; | ||
|
|
||
| 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(); | ||
| } | ||
|
|
||
| quark_skein512_cpu_init(thr_id, throughput); | ||
| x11_shavite512_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; | ||
| } | ||
|
|
||
| uint32_t endiandata[20]; | ||
| for (int k=0; k < 20; k++) | ||
| be32enc(&endiandata[k], pdata[k]); | ||
|
|
||
| skein512_cpu_setBlock_80(endiandata); | ||
| cuda_check_cpu_setTarget(ptarget); | ||
|
|
||
| do { | ||
| int order = 0; | ||
| uint32_t foundNonce; | ||
|
|
||
| // Hash with CUDA | ||
| skein512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], 1); order++; | ||
| TRACE("blake :"); | ||
| x11_shavite512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
| TRACE("shavite:"); | ||
| x14_shabal512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
| TRACE("shabal :"); | ||
| streebog_cpu_hash_64(thr_id, throughput, pdata[19], d_hash[thr_id]); | ||
| TRACE("gost :"); | ||
|
|
||
| *hashes_done = pdata[19] - first_nonce + throughput; | ||
|
|
||
| foundNonce = cuda_check_hash(thr_id, throughput, pdata[19], d_hash[thr_id]); | ||
| if (foundNonce != UINT32_MAX) | ||
| { | ||
| const uint32_t Htarg = ptarget[7]; | ||
| uint32_t _ALIGN(64) vhash[8]; | ||
| be32enc(&endiandata[19], foundNonce); | ||
| veltorhash(vhash, endiandata); | ||
|
|
||
| if (vhash[7] <= Htarg && fulltest(vhash, ptarget)) { | ||
| int res = 1; | ||
| // check if there was another one... | ||
| uint32_t secNonce = cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], 1); | ||
| work_set_target_ratio(work, vhash); | ||
| if (secNonce != 0) { | ||
| be32enc(&endiandata[19], secNonce); | ||
| veltorhash(vhash, endiandata); | ||
| work->nonces[1] = secNonce; | ||
| if (bn_hash_target_ratio(vhash, ptarget) > work->shareratio) { | ||
| work_set_target_ratio(work, vhash); | ||
| xchg(work->nonces[1], work->nonces[0]); | ||
| } | ||
| res++; | ||
| } | ||
| pdata[19] = work->nonces[0] = foundNonce; | ||
| return res; | ||
| } else { | ||
| gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce); | ||
| pdata[19] = foundNonce + 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_veltor(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(); | ||
| } |