int scanhash_axiom(int thr_id, struct work *work, uint32_t max_nonce, uint64_t *hashes_done) { uint32_t *pdata = work->data; uint32_t *ptarget = work->target; uint32_t _ALIGN(64) hash64[8]; uint32_t _ALIGN(64) endiandata[20]; const uint32_t Htarg = ptarget[7]; const uint32_t first_nonce = pdata[19]; uint32_t n = first_nonce; for (int k = 0; k < 19; k++) be32enc(&endiandata[k], pdata[k]); do { be32enc(&endiandata[19], n); axiomhash(hash64, endiandata); if (hash64[7] < Htarg && fulltest(hash64, ptarget)) { *hashes_done = n - first_nonce + 1; pdata[19] = n; work_set_target_ratio( work, hash64 ); return true; } n++; } while (n < max_nonce && !work_restart[thr_id].restart); *hashes_done = n - first_nonce + 1; pdata[19] = n; return 0; }
int scanhash_drop(int thr_id, struct work *work, uint32_t max_nonce, uint64_t *hashes_done) { uint32_t _ALIGN(128) hash[16]; uint32_t *pdata = work->data; uint32_t *ptarget = work->target; const uint32_t version = pdata[0] & (~POK_DATA_MASK); const uint32_t first_nonce = pdata[19]; uint32_t nonce = first_nonce; #define tmpdata pdata if (opt_benchmark) ptarget[7] = 0x07ff; const uint32_t htarg = ptarget[7]; do { tmpdata[19] = nonce; droplp_hash_pok(hash, tmpdata, version); if (hash[7] <= htarg && fulltest(hash, ptarget)) { work_set_target_ratio(work, hash); pdata[0] = tmpdata[0]; pdata[19] = nonce; *hashes_done = pdata[19] - first_nonce + 1; return 1; } nonce++; } while (nonce < max_nonce && !work_restart[thr_id].restart); pdata[19] = nonce; *hashes_done = pdata[19] - first_nonce + 1; return 0; }
int scanhash_zoin( int thr_id, struct work *work, uint32_t max_nonce, uint64_t *hashes_done ) { uint32_t _ALIGN(128) hash[8]; uint32_t _ALIGN(128) endiandata[20]; uint32_t *pdata = work->data; uint32_t *ptarget = work->target; const uint32_t Htarg = ptarget[7]; const uint32_t first_nonce = pdata[19]; uint32_t nonce = first_nonce; if (opt_benchmark) ptarget[7] = 0x0000ff; for (int i=0; i < 19; i++) { be32enc(&endiandata[i], pdata[i]); } do { be32enc(&endiandata[19], nonce); zoin_hash( hash, endiandata, work->height ); if (hash[7] <= Htarg && fulltest(hash, ptarget)) { work_set_target_ratio(work, hash); pdata[19] = nonce; *hashes_done = pdata[19] - first_nonce; return 1; } nonce++; } while (nonce < max_nonce && !work_restart[thr_id].restart); pdata[19] = nonce; *hashes_done = pdata[19] - first_nonce + 1; return 0; }
int scanhash_x12_4way( int thr_id, struct work *work, uint32_t max_nonce, uint64_t *hashes_done ) { uint32_t hash[4*8] __attribute__ ((aligned (64))); uint32_t vdata[24*4] __attribute__ ((aligned (64))); uint32_t endiandata[20] __attribute__((aligned(64))); uint32_t *pdata = work->data; uint32_t *ptarget = work->target; uint32_t n = pdata[19]; const uint32_t first_nonce = pdata[19]; uint32_t *nonces = work->nonces; int num_found = 0; uint32_t *noncep = vdata + 73; // 9*8 + 1 const uint32_t Htarg = ptarget[7]; uint64_t htmax[] = { 0, 0xF, 0xFF, 0xFFF, 0xFFFF, 0x10000000 }; uint32_t masks[] = { 0xFFFFFFFF, 0xFFFFFFF0, 0xFFFFFF00, 0xFFFFF000, 0xFFFF0000, 0 }; // big endian encode 0..18 uint32_t, 64 bits at a time swab32_array( endiandata, pdata, 20 ); uint64_t *edata = (uint64_t*)endiandata; mm256_interleave_4x64( (uint64_t*)vdata, edata, edata, edata, edata, 640 ); for ( int m=0; m < 6; m++ ) if ( Htarg <= htmax[m] ) { uint32_t mask = masks[m]; do { be32enc( noncep, n ); be32enc( noncep+2, n+1 ); be32enc( noncep+4, n+2 ); be32enc( noncep+6, n+3 ); x12_4way_hash( hash, vdata ); pdata[19] = n; for ( int i = 0; i < 4; i++ ) if ( ( ( (hash+(i<<3))[7] & mask ) == 0 ) && fulltest( hash+(i<<3), ptarget ) ) { pdata[19] = n+i; nonces[ num_found++ ] = n+i; work_set_target_ratio( work, hash+(i<<3) ); } n += 4; } while ( ( num_found == 0 ) && ( n < max_nonce ) && !work_restart[thr_id].restart ); break; } *hashes_done = n - first_nonce + 1; return num_found; }
int scanhash_decred(int thr_id, struct work *work, uint32_t max_nonce, uint64_t *hashes_done) { uint32_t _ALIGN(128) endiandata[48]; uint32_t _ALIGN(128) hash32[8]; uint32_t *pdata = work->data; uint32_t *ptarget = work->target; #define DCR_NONCE_OFT32 35 const uint32_t first_nonce = pdata[DCR_NONCE_OFT32]; const uint32_t HTarget = opt_benchmark ? 0x7f : ptarget[7]; uint32_t n = first_nonce; ctx_midstate_done = false; #if 1 memcpy(endiandata, pdata, 180); #else for (int k=0; k < (180/4); k++) be32enc(&endiandata[k], pdata[k]); #endif #ifdef DEBUG_ALGO if (!thr_id) applog(LOG_DEBUG,"[%d] Target=%08x %08x", thr_id, ptarget[6], ptarget[7]); #endif do { //be32enc(&endiandata[DCR_NONCE_OFT32], n); endiandata[DCR_NONCE_OFT32] = n; decred_hash(hash32, endiandata); if (hash32[7] <= HTarget && fulltest(hash32, ptarget)) { work_set_target_ratio(work, hash32); *hashes_done = n - first_nonce + 1; #ifdef DEBUG_ALGO applog(LOG_BLUE, "Nonce : %08x %08x", n, swab32(n)); applog_hash(ptarget); applog_compare_hash(hash32, ptarget); #endif pdata[DCR_NONCE_OFT32] = n; return 1; } n++; } while (n < max_nonce && !work_restart[thr_id].restart); *hashes_done = n - first_nonce + 1; pdata[DCR_NONCE_OFT32] = n; return 0; }
int scanhash_veltor_4way( int thr_id, struct work *work, uint32_t max_nonce, uint64_t *hashes_done ) { uint32_t hash[4*8] __attribute__ ((aligned (64))); uint32_t vdata[24*4] __attribute__ ((aligned (64))); uint32_t endiandata[20] __attribute__((aligned(64))); uint32_t *pdata = work->data; uint32_t *ptarget = work->target; const uint32_t Htarg = ptarget[7]; const uint32_t first_nonce = pdata[19]; uint32_t n = first_nonce; uint32_t *nonces = work->nonces; int num_found = 0; uint32_t *noncep = vdata + 73; // 9*8 + 1 volatile uint8_t *restart = &(work_restart[thr_id].restart); if ( opt_benchmark ) ptarget[7] = 0x0cff; for ( int i=0; i < 19; i++ ) { be32enc( &endiandata[i], pdata[i] ); } uint64_t *edata = (uint64_t*)endiandata; mm256_interleave_4x64( (uint64_t*)vdata, edata, edata, edata, edata, 640 ); do { be32enc( noncep, n ); be32enc( noncep+2, n+1 ); be32enc( noncep+4, n+2 ); be32enc( noncep+6, n+3 ); veltor_4way_hash( hash, vdata ); pdata[19] = n; for ( int i = 0; i < 4; i++ ) if ( (hash+(i<<3))[7] <= Htarg && fulltest( hash+(i<<3), ptarget ) ) { pdata[19] = n+i; nonces[ num_found++ ] = n+i; work_set_target_ratio( work, hash+(i<<3) ); } n += 4; } while ( ( num_found == 0 ) && ( n < max_nonce ) && !(*restart) ); *hashes_done = n - first_nonce + 1; return num_found; }
int scanhash_lyra2z(int thr_id, struct work *work, uint32_t max_nonce, uint64_t *hashes_done) { size_t size = (int64_t) ((int64_t) 16 * 16 * 96); uint64_t *wholeMatrix = _mm_malloc(size, 64); uint32_t _ALIGN(128) hash[8]; uint32_t _ALIGN(128) endiandata[20]; uint32_t *pdata = work->data; uint32_t *ptarget = work->target; const uint32_t Htarg = ptarget[7]; const uint32_t first_nonce = pdata[19]; uint32_t nonce = first_nonce; if (opt_benchmark) ptarget[7] = 0x0000ff; for (int i=0; i < 19; i++) { be32enc(&endiandata[i], pdata[i]); } do { be32enc(&endiandata[19], nonce); lyra2z_hash(wholeMatrix, hash, endiandata); // lyra2z_hash(0, hash, endiandata); if (hash[7] <= Htarg && fulltest(hash, ptarget)) { work_set_target_ratio(work, hash); pdata[19] = nonce; *hashes_done = pdata[19] - first_nonce; _mm_free(wholeMatrix); return 1; } nonce++; } while (nonce < max_nonce && !work_restart[thr_id].restart); pdata[19] = nonce; *hashes_done = pdata[19] - first_nonce + 1; _mm_free(wholeMatrix); return 0; }
int scanhash_x16r(int thr_id, struct work *work, uint32_t max_nonce, uint64_t *hashes_done) { uint32_t _ALIGN(128) hash32[8]; uint32_t _ALIGN(128) endiandata[20]; uint32_t *pdata = work->data; uint32_t *ptarget = work->target; const uint32_t Htarg = ptarget[7]; const uint32_t first_nonce = pdata[19]; uint32_t nonce = first_nonce; volatile uint8_t *restart = &(work_restart[thr_id].restart); for (int k=0; k < 19; k++) be32enc(&endiandata[k], pdata[k]); if (s_ntime != pdata[17]) { uint32_t ntime = swab32(pdata[17]); getAlgoString((const char*) (&endiandata[1]), hashOrder); s_ntime = ntime; if (opt_debug && !thr_id) applog(LOG_DEBUG, "hash order %s (%08x)", hashOrder, ntime); } if (opt_benchmark) ptarget[7] = 0x0cff; do { be32enc(&endiandata[19], nonce); x16r_hash(hash32, endiandata); if (hash32[7] <= Htarg && fulltest(hash32, ptarget)) { work_set_target_ratio(work, hash32); pdata[19] = nonce; *hashes_done = pdata[19] - first_nonce; return 1; } nonce++; } while (nonce < max_nonce && !(*restart)); pdata[19] = nonce; *hashes_done = pdata[19] - first_nonce + 1; return 0; }
int scanhash_blake(int thr_id, struct work *work, uint32_t max_nonce, uint64_t *hashes_done) { uint32_t _ALIGN(128) hash32[8]; uint32_t _ALIGN(128) endiandata[20]; uint32_t *pdata = work->data; uint32_t *ptarget = work->target; const uint32_t first_nonce = pdata[19]; const uint32_t HTarget = opt_benchmark ? 0x7f : ptarget[7]; uint32_t n = first_nonce; ctx_midstate_done = false; // we need big endian data... for (int kk=0; kk < 19; kk++) { be32enc(&endiandata[kk], pdata[kk]); } #ifdef DEBUG_ALGO applog(LOG_DEBUG,"[%d] Target=%08x %08x", thr_id, ptarget[6], ptarget[7]); #endif do { be32enc(&endiandata[19], n); blakehash(hash32, endiandata); if (hash32[7] <= HTarget && fulltest(hash32, ptarget)) { work_set_target_ratio(work, hash32); *hashes_done = n - first_nonce + 1; return 1; } n++; pdata[19] = n; } while (n < max_nonce && !work_restart[thr_id].restart); *hashes_done = n - first_nonce + 1; pdata[19] = n; return 0; }
int scanhash_qubit(int thr_id, struct work *work, uint32_t max_nonce, uint64_t *hashes_done) { uint32_t _ALIGN(128) hash32[8]; uint32_t _ALIGN(128) endiandata[20]; uint32_t *pdata = work->data; uint32_t *ptarget = work->target; uint32_t n = pdata[19] - 1; const uint32_t first_nonce = pdata[19]; const uint32_t Htarg = ptarget[7]; uint64_t htmax[] = { 0, 0xF, 0xFF, 0xFFF, 0xFFFF, 0x10000000 }; uint32_t masks[] = { 0xFFFFFFFF, 0xFFFFFFF0, 0xFFFFFF00, 0xFFFFF000, 0xFFFF0000, 0 }; // we need bigendian data... for (int i=0; i < 19; i++) { be32enc(&endiandata[i], pdata[i]); } #ifdef DEBUG_ALGO printf("[%d] Htarg=%X\n", thr_id, Htarg); #endif for (int m=0; m < 6; m++) { if (Htarg <= htmax[m]) { uint32_t mask = masks[m]; do { pdata[19] = ++n; be32enc(&endiandata[19], n); qubithash(hash32, endiandata); #ifndef DEBUG_ALGO if ((!(hash32[7] & mask)) && fulltest(hash32, ptarget)) { work_set_target_ratio(work, hash32); *hashes_done = n - first_nonce + 1; return 1; } #else if (!(n % 0x1000) && !thr_id) printf("."); if (!(hash32[7] & mask)) { printf("[%d]",thr_id); if (fulltest(hash32, ptarget)) { work_set_target_ratio(work, hash32); *hashes_done = n - first_nonce + 1; return 1; } } #endif } while (n < max_nonce && !work_restart[thr_id].restart); // see blake.c if else to understand the loop on htmax => mask break; } } *hashes_done = n - first_nonce + 1; pdata[19] = n; return 0; }
int scanhash_myriad(int thr_id, struct work *work, uint32_t max_nonce, unsigned long *hashes_done) { uint32_t _ALIGN(64) endiandata[32]; uint32_t *pdata = work->data; uint32_t *ptarget = work->target; uint32_t start_nonce = pdata[19]; uint32_t throughput = cuda_default_throughput(thr_id, 1U << 17); if (init[thr_id]) throughput = min(throughput, max_nonce - start_nonce); uint32_t *outputHash = (uint32_t*)malloc(throughput * 64); if (opt_benchmark) ptarget[7] = 0x0000ff; // init 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(); } myriadgroestl_cpu_init(thr_id, throughput); init[thr_id] = true; } for (int k=0; k < 20; k++) be32enc(&endiandata[k], pdata[k]); // Context mit dem Endian gedrehten Blockheader vorbereiten (Nonce wird später ersetzt) myriadgroestl_cpu_setBlock(thr_id, endiandata, (void*)ptarget); do { // GPU uint32_t foundNounce = UINT32_MAX; myriadgroestl_cpu_hash(thr_id, throughput, pdata[19], outputHash, &foundNounce); *hashes_done = pdata[19] - start_nonce + throughput; if (foundNounce < UINT32_MAX && bench_algo < 0) { uint32_t _ALIGN(64) vhash[8]; endiandata[19] = swab32(foundNounce); myriadhash(vhash, endiandata); if (vhash[7] <= ptarget[7] && fulltest(vhash, ptarget)) { work_set_target_ratio(work, vhash); pdata[19] = foundNounce; free(outputHash); return 1; } else { gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNounce); } } if ((uint64_t) throughput + pdata[19] >= max_nonce) { pdata[19] = max_nonce; break; } pdata[19] += throughput; } while (!work_restart[thr_id].restart); *hashes_done = max_nonce - start_nonce; free(outputHash); return 0; }
int scanhash_skein2(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done) { int dev_id = device_map[thr_id]; uint32_t *pdata = work->data; uint32_t *ptarget = work->target; const uint32_t first_nonce = pdata[19]; const int swap = 1; // to toggle nonce endian uint32_t throughput = cuda_default_throughput(thr_id, 1U << 19); // 256*256*8 if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce); if (opt_benchmark) ((uint32_t*)ptarget)[7] = 0; if (!init[thr_id]) { cudaSetDevice(dev_id); if (opt_cudaschedule == -1 && gpu_threads == 1) { cudaDeviceReset(); // reduce cpu usage cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); CUDA_LOG_ERROR(); } cudaMalloc(&d_hash[thr_id], (size_t) 64 * throughput); quark_skein512_cpu_init(thr_id, throughput); cuda_check_cpu_init(thr_id, throughput); CUDA_SAFE_CALL(cudaDeviceSynchronize()); init[thr_id] = true; } uint32_t endiandata[20]; for (int k=0; k < 19; k++) be32enc(&endiandata[k], pdata[k]); skein512_cpu_setBlock_80((void*)endiandata); cuda_check_cpu_setTarget(ptarget); do { int order = 0; // Hash with CUDA skein512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], swap); quark_skein512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); *hashes_done = pdata[19] - first_nonce + throughput; uint32_t foundNonce = cuda_check_hash(thr_id, throughput, pdata[19], d_hash[thr_id]); if (foundNonce != UINT32_MAX) { uint32_t _ALIGN(64) vhash64[8]; endiandata[19] = swab32_if(foundNonce, swap); skein2hash(vhash64, endiandata); if (vhash64[7] <= ptarget[7] && fulltest(vhash64, ptarget)) { int res = 1; uint32_t secNonce = cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], 1); work_set_target_ratio(work, vhash64); if (secNonce != 0) { if (!opt_quiet) applog(LOG_BLUE, "GPU #%d: found second nonce %08x !", dev_id, swab32(secNonce)); endiandata[19] = swab32_if(secNonce, swap); skein2hash(vhash64, endiandata); if (bn_hash_target_ratio(vhash64, ptarget) > work->shareratio) work_set_target_ratio(work, vhash64); pdata[21] = swab32_if(secNonce, !swap); res++; } pdata[19] = swab32_if(foundNonce, !swap); return res; } else { gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce); } } 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; }
int scanhash_fugue256(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; uint32_t start_nonce = pdata[19]++; int intensity = (device_sm[device_map[thr_id]] > 500) ? 22 : 19; uint32_t throughput = cuda_default_throughput(thr_id, 1U << intensity); if (init[thr_id]) throughput = min(throughput, max_nonce - start_nonce); if (opt_benchmark) ptarget[7] = 0xf; // init if(!init[thr_id]) { cudaSetDevice(device_map[thr_id]); fugue256_cpu_init(thr_id, throughput); init[thr_id] = true; } // Endian for (int kk=0; kk < 20; kk++) be32enc(&endiandata[kk], pdata[kk]); fugue256_cpu_setBlock(thr_id, endiandata, (void*)ptarget); do { // GPU uint32_t foundNounce = UINT32_MAX; fugue256_cpu_hash(thr_id, throughput, pdata[19], NULL, &foundNounce); *hashes_done = pdata[19] - start_nonce + throughput; if (foundNounce < UINT32_MAX && bench_algo < 0) { uint32_t vhash[8]; sph_fugue256_context ctx_fugue; endiandata[19] = SWAP32(foundNounce); sph_fugue256_init(&ctx_fugue); sph_fugue256 (&ctx_fugue, endiandata, 80); sph_fugue256_close(&ctx_fugue, &vhash); if (vhash[7] <= ptarget[7] && fulltest(vhash, ptarget)) { work_set_target_ratio(work, vhash); pdata[19] = foundNounce; return 1; } else { gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNounce); } } 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] - start_nonce; return 0; }