int scanhash_skein2(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; swab32_array( endiandata, pdata, 20 ); do { be32enc(&endiandata[19], n); skein2hash(hash64, endiandata); if (hash64[7] < Htarg && fulltest(hash64, ptarget)) { *hashes_done = n - first_nonce + 1; pdata[19] = n; 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_skein2(int thr_id, uint32_t *pdata, const uint32_t *ptarget, uint32_t max_nonce, uint64_t *hashes_done) { 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 i=0; i < 19; i++) { be32enc(&endiandata[i], pdata[i]); }; do { be32enc(&endiandata[19], n); skein2hash(hash64, endiandata); if (hash64[7] < Htarg && fulltest(hash64, ptarget)) { *hashes_done = n - first_nonce + 1; pdata[19] = n; 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_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; }