Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
committing a quick attempt at NIST5 (TalkCoin)
- Loading branch information
Showing
6 changed files
with
194 additions
and
11 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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,161 @@ | ||
|
|
||
| extern "C" | ||
| { | ||
| #include "sph/sph_blake.h" | ||
| #include "sph/sph_groestl.h" | ||
| #include "sph/sph_skein.h" | ||
| #include "sph/sph_jh.h" | ||
| #include "sph/sph_keccak.h" | ||
| #include "miner.h" | ||
| } | ||
|
|
||
| #include <stdint.h> | ||
|
|
||
| // aus cpu-miner.c | ||
| extern int device_map[8]; | ||
|
|
||
| // Speicher für Input/Output der verketteten Hashfunktionen | ||
| static uint32_t *d_hash[8]; | ||
|
|
||
| extern void quark_blake512_cpu_init(int thr_id, int threads); | ||
| extern void quark_blake512_cpu_setBlock_80(void *pdata); | ||
| extern void quark_blake512_cpu_hash_80(int thr_id, int threads, uint32_t startNounce, uint32_t *d_hash, int order); | ||
|
|
||
| extern void quark_groestl512_cpu_init(int thr_id, int threads); | ||
| extern void quark_groestl512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); | ||
|
|
||
| extern void quark_jh512_cpu_init(int thr_id, int threads); | ||
| extern void quark_jh512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); | ||
|
|
||
| extern void quark_keccak512_cpu_init(int thr_id, int threads); | ||
| extern void quark_keccak512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); | ||
|
|
||
| extern void quark_skein512_cpu_init(int thr_id, int threads); | ||
| extern void quark_skein512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); | ||
|
|
||
| extern void quark_check_cpu_init(int thr_id, int threads); | ||
| extern void quark_check_cpu_setTarget(const void *ptarget); | ||
| extern uint32_t quark_check_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_inputHash, int order); | ||
|
|
||
| // Original nist5hash Funktion aus einem miner Quelltext | ||
| inline void nist5hash(void *state, const void *input) | ||
| { | ||
| sph_blake512_context ctx_blake; | ||
| sph_groestl512_context ctx_groestl; | ||
| sph_jh512_context ctx_jh; | ||
| sph_keccak512_context ctx_keccak; | ||
| sph_skein512_context ctx_skein; | ||
|
|
||
| unsigned char hash[64]; | ||
|
|
||
| sph_blake512_init(&ctx_blake); | ||
| // ZBLAKE; | ||
| sph_blake512 (&ctx_blake, input, 80); | ||
| sph_blake512_close(&ctx_blake, (void*) hash); | ||
|
|
||
| sph_groestl512_init(&ctx_groestl); | ||
| // ZGROESTL; | ||
| sph_groestl512 (&ctx_groestl, (const void*) hash, 64); | ||
| sph_groestl512_close(&ctx_groestl, (void*) hash); | ||
|
|
||
| sph_jh512_init(&ctx_jh); | ||
| // ZJH; | ||
| sph_jh512 (&ctx_jh, (const void*) hash, 64); | ||
| sph_jh512_close(&ctx_jh, (void*) hash); | ||
|
|
||
| sph_keccak512_init(&ctx_keccak); | ||
| // ZKECCAK; | ||
| sph_keccak512 (&ctx_keccak, (const void*) hash, 64); | ||
| sph_keccak512_close(&ctx_keccak, (void*) hash); | ||
|
|
||
| sph_skein512_init(&ctx_skein); | ||
| // ZSKEIN; | ||
| sph_skein512 (&ctx_skein, (const void*) hash, 64); | ||
| sph_skein512_close(&ctx_skein, (void*) hash); | ||
|
|
||
| memcpy(state, hash, 32); | ||
| } | ||
|
|
||
|
|
||
| extern bool opt_benchmark; | ||
|
|
||
| extern "C" int scanhash_nist5(int thr_id, uint32_t *pdata, | ||
| const uint32_t *ptarget, uint32_t max_nonce, | ||
| unsigned long *hashes_done) | ||
| { | ||
| const uint32_t first_nonce = pdata[19]; | ||
|
|
||
| // TODO: entfernen für eine Release! Ist nur zum Testen! | ||
| if (opt_benchmark) | ||
| ((uint32_t*)ptarget)[7] = 0x0000ff; | ||
|
|
||
| const uint32_t Htarg = ptarget[7]; | ||
|
|
||
| const int throughput = 256*4096; // 100; | ||
|
|
||
| static bool init[8] = {0,0,0,0,0,0,0,0}; | ||
| if (!init[thr_id]) | ||
| { | ||
| cudaSetDevice(device_map[thr_id]); | ||
|
|
||
| // Konstanten kopieren, Speicher belegen | ||
| cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput); | ||
| quark_blake512_cpu_init(thr_id, throughput); | ||
| quark_groestl512_cpu_init(thr_id, throughput); | ||
| quark_jh512_cpu_init(thr_id, throughput); | ||
| quark_keccak512_cpu_init(thr_id, throughput); | ||
| quark_skein512_cpu_init(thr_id, throughput); | ||
| quark_check_cpu_init(thr_id, throughput); | ||
| init[thr_id] = true; | ||
| } | ||
|
|
||
| uint32_t endiandata[20]; | ||
| for (int k=0; k < 20; k++) | ||
| be32enc(&endiandata[k], ((uint32_t*)pdata)[k]); | ||
|
|
||
| quark_blake512_cpu_setBlock_80((void*)endiandata); | ||
| quark_check_cpu_setTarget(ptarget); | ||
|
|
||
| do { | ||
| int order = 0; | ||
|
|
||
| // erstes Blake512 Hash mit CUDA | ||
| quark_blake512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++); | ||
|
|
||
| // das ist der unbedingte Branch für Groestl512 | ||
| quark_groestl512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
|
|
||
| // das ist der unbedingte Branch für JH512 | ||
| quark_jh512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
|
|
||
| // das ist der unbedingte Branch für Keccak512 | ||
| quark_keccak512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
|
|
||
| // das ist der unbedingte Branch für Skein512 | ||
| quark_skein512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
|
|
||
| // Scan nach Gewinner Hashes auf der GPU | ||
| uint32_t foundNonce = quark_check_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); | ||
| if (foundNonce != 0xffffffff) | ||
| { | ||
| uint32_t vhash64[8]; | ||
| be32enc(&endiandata[19], foundNonce); | ||
| nist5hash(vhash64, endiandata); | ||
|
|
||
| if ((vhash64[7]<=Htarg) && fulltest(vhash64, ptarget)) { | ||
|
|
||
| pdata[19] = foundNonce; | ||
| *hashes_done = foundNonce - first_nonce + 1; | ||
| return 1; | ||
| } else { | ||
| applog(LOG_INFO, "GPU #%d: result for nonce $%08X does not validate on CPU!", thr_id, foundNonce); | ||
| } | ||
| } | ||
|
|
||
| pdata[19] += throughput; | ||
|
|
||
| } while (pdata[19] < max_nonce && !work_restart[thr_id].restart); | ||
|
|
||
| *hashes_done = pdata[19] - first_nonce + 1; | ||
| return 0; | ||
| } |
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