Browse Source

never interrupt global benchmark with found nonces

fix some algo weird hashrates (like blake)
and reset device between algos, for better accuracy

but this reset doesnt seems enough to bench all algos correctly...

to test on linux, could be a driver issue...

heavy: fix first alloc and indent with tabs...
2upstream
Tanguy Pruvot 9 years ago
parent
commit
61ff92b5b4
  1. 5
      Algo256/blake256.cu
  2. 2
      Algo256/bmw.cu
  3. 8
      Algo256/keccak256.cu
  4. 7
      JHA/jackpotcoin.cu
  5. 2
      Makefile.am
  6. 10
      bench.cpp
  7. 15
      ccminer.cpp
  8. 4
      cuda.cpp
  9. 11
      cuda_checkhash.cu
  10. 8
      cuda_nist5.cu
  11. 9
      fuguecoin.cpp
  12. 11
      groestlcoin.cpp
  13. 575
      heavy/heavy.cu
  14. 11
      lyra2/lyra2RE.cu
  15. 11
      lyra2/lyra2REv2.cu
  16. 13
      myriadgroestl.cpp
  17. 9
      neoscrypt/neoscrypt.cpp
  18. 7
      pentablake.cu
  19. 7
      qubit/deep.cu
  20. 25
      qubit/luffa.cu
  21. 6
      skein.cu
  22. 5
      skein2.cpp
  23. 12
      x11/c11.cu
  24. 11
      x11/fresh.cu
  25. 9
      x11/s3.cu
  26. 11
      x11/x11.cu
  27. 4
      x13/x13.cu
  28. 11
      x15/whirlpool.cu
  29. 20
      x15/whirlpoolx.cu
  30. 10
      x15/x14.cu
  31. 9
      x15/x15.cu
  32. 7
      x17/x17.cu

5
Algo256/blake256.cu

@ -439,7 +439,7 @@ extern "C" int scanhash_blake256(int thr_id, struct work* work, uint32_t max_non
#endif #endif
*hashes_done = pdata[19] - first_nonce + throughput; *hashes_done = pdata[19] - first_nonce + throughput;
if (foundNonce != UINT32_MAX) if (foundNonce != UINT32_MAX && bench_algo == -1)
{ {
uint32_t vhashcpu[8]; uint32_t vhashcpu[8];
uint32_t Htarg = (uint32_t)targetHigh; uint32_t Htarg = (uint32_t)targetHigh;
@ -478,7 +478,8 @@ extern "C" int scanhash_blake256(int thr_id, struct work* work, uint32_t max_non
} }
} }
if ((uint64_t) pdata[19] + throughput > (uint64_t) max_nonce) { if ((uint64_t) throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
break; break;
} }

2
Algo256/bmw.cu

@ -91,7 +91,7 @@ extern "C" int scanhash_bmw(int thr_id, struct work* work, uint32_t max_nonce, u
} }
} }
if ((uint64_t) throughput + pdata[19] > max_nonce) { if ((uint64_t) throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce; pdata[19] = max_nonce;
break; break;
} }

8
Algo256/keccak256.cu

@ -46,7 +46,7 @@ extern "C" int scanhash_keccak256(int thr_id, struct work* work, uint32_t max_no
if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce); if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);
if (opt_benchmark) if (opt_benchmark)
ptarget[7] = 0x00ff; ptarget[7] = 0x000f;
if (!init[thr_id]) { if (!init[thr_id]) {
cudaSetDevice(device_map[thr_id]); cudaSetDevice(device_map[thr_id]);
@ -68,7 +68,7 @@ extern "C" int scanhash_keccak256(int thr_id, struct work* work, uint32_t max_no
*hashes_done = pdata[19] - first_nonce + throughput; *hashes_done = pdata[19] - first_nonce + throughput;
uint32_t foundNonce = keccak256_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++); uint32_t foundNonce = keccak256_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
if (foundNonce != UINT32_MAX) if (foundNonce != UINT32_MAX && bench_algo < 0)
{ {
uint32_t _ALIGN(64) vhash64[8]; uint32_t _ALIGN(64) vhash64[8];
be32enc(&endiandata[19], foundNonce); be32enc(&endiandata[19], foundNonce);
@ -84,7 +84,8 @@ extern "C" int scanhash_keccak256(int thr_id, struct work* work, uint32_t max_no
} }
} }
if ((uint64_t) pdata[19] + throughput > max_nonce) { if ((uint64_t) throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
break; break;
} }
@ -92,6 +93,7 @@ extern "C" int scanhash_keccak256(int thr_id, struct work* work, uint32_t max_no
} while (!work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce;
return 0; return 0;
} }

7
JHA/jackpotcoin.cu

@ -238,16 +238,17 @@ extern "C" int scanhash_jackpot(int thr_id, struct work *work, uint32_t max_nonc
} }
} }
if ((uint64_t) pdata[19] + throughput > max_nonce) { if ((uint64_t) throughput + pdata[19] >= max_nonce) {
*hashes_done = pdata[19] - first_nonce;
pdata[19] = max_nonce; pdata[19] = max_nonce;
return 0; break;
} }
pdata[19] += throughput; pdata[19] += throughput;
} while (!work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce;
CUDA_LOG_ERROR(); CUDA_LOG_ERROR();
return 0; return 0;

2
Makefile.am

@ -81,7 +81,7 @@ ccminer_CPPFLAGS = @LIBCURL_CPPFLAGS@ @PCIFLAGS@ @OPENMP_CFLAGS@ $(CPPFLAGS) $(P
nvcc_ARCH = -gencode=arch=compute_50,code=\"sm_50,compute_50\" nvcc_ARCH = -gencode=arch=compute_50,code=\"sm_50,compute_50\"
#nvcc_ARCH += -gencode=arch=compute_52,code=\"sm_52,compute_52\" nvcc_ARCH += -gencode=arch=compute_52,code=\"sm_52,compute_52\"
#nvcc_ARCH += -gencode=arch=compute_35,code=\"sm_35,compute_35\" #nvcc_ARCH += -gencode=arch=compute_35,code=\"sm_35,compute_35\"
#nvcc_ARCH += -gencode=arch=compute_30,code=\"sm_30,compute_30\" #nvcc_ARCH += -gencode=arch=compute_30,code=\"sm_30,compute_30\"
#nvcc_ARCH += -gencode=arch=compute_20,code=\"sm_21,compute_20\" #nvcc_ARCH += -gencode=arch=compute_20,code=\"sm_21,compute_20\"

10
bench.cpp

@ -89,6 +89,9 @@ bool bench_algo_switch_next(int thr_id)
int prev_algo = algo; int prev_algo = algo;
int dev_id = device_map[thr_id % MAX_GPUS]; int dev_id = device_map[thr_id % MAX_GPUS];
int mfree, mused; int mfree, mused;
// doesnt seems enough to prevent device slow down
// after some algo switchs
bool need_reset = (gpu_threads == 1);
algo++; algo++;
@ -143,6 +146,7 @@ bool bench_algo_switch_next(int thr_id)
gpulog(LOG_WARNING, thr_id, "possible %d MB memory leak in %s! %d MB free", gpulog(LOG_WARNING, thr_id, "possible %d MB memory leak in %s! %d MB free",
(device_mem_free[thr_id] - mfree), algo_names[prev_algo], mfree); (device_mem_free[thr_id] - mfree), algo_names[prev_algo], mfree);
cuda_reset_device(thr_id, NULL); // force to free the leak cuda_reset_device(thr_id, NULL); // force to free the leak
need_reset = false;
mfree = cuda_available_memory(thr_id); mfree = cuda_available_memory(thr_id);
} }
// store used memory per algo // store used memory per algo
@ -152,14 +156,13 @@ bool bench_algo_switch_next(int thr_id)
// store to dump a table per gpu later // store to dump a table per gpu later
algo_hashrates[thr_id][prev_algo] = hashrate; algo_hashrates[thr_id][prev_algo] = hashrate;
// wait the other threads to display logs correctly // wait the other threads to display logs correctly
if (opt_n_threads > 1) { if (opt_n_threads > 1) {
pthread_barrier_wait(&algo_barr); pthread_barrier_wait(&algo_barr);
} }
if (algo == ALGO_AUTO) if (algo == ALGO_AUTO)
return false; return false; // all algos done
// mutex primary used for the stats purge // mutex primary used for the stats purge
pthread_mutex_lock(&bench_lock); pthread_mutex_lock(&bench_lock);
@ -170,6 +173,9 @@ bool bench_algo_switch_next(int thr_id)
thr_hashrates[thr_id] = 0; // reset for minmax64 thr_hashrates[thr_id] = 0; // reset for minmax64
pthread_mutex_unlock(&bench_lock); pthread_mutex_unlock(&bench_lock);
if (need_reset)
cuda_reset_device(thr_id, NULL);
if (thr_id == 0) if (thr_id == 0)
applog(LOG_BLUE, "Benchmark algo %s...", algo_names[algo]); applog(LOG_BLUE, "Benchmark algo %s...", algo_names[algo]);

15
ccminer.cpp

@ -479,9 +479,11 @@ void proper_exit(int reason)
reason = app_exit_code; reason = app_exit_code;
} }
pthread_mutex_lock(&stats_lock);
if (check_dups) if (check_dups)
hashlog_purge_all(); hashlog_purge_all();
stats_purge_all(); stats_purge_all();
pthread_mutex_unlock(&stats_lock);
#ifdef WIN32 #ifdef WIN32
timeEndPeriod(1); // else never executed timeEndPeriod(1); // else never executed
@ -496,7 +498,7 @@ void proper_exit(int reason)
#endif #endif
free(opt_syslog_pfx); free(opt_syslog_pfx);
free(opt_api_allow); free(opt_api_allow);
free(work_restart); //free(work_restart);
//free(thr_info); //free(thr_info);
exit(reason); exit(reason);
} }
@ -1709,18 +1711,22 @@ static void *miner_thread(void *userdata)
if (max64 < minmax) { if (max64 < minmax) {
switch (opt_algo) { switch (opt_algo) {
case ALGO_BLAKECOIN: case ALGO_BLAKECOIN:
case ALGO_BLAKE:
case ALGO_WHIRLPOOLX:
minmax = 0x80000000U; minmax = 0x80000000U;
break; break;
case ALGO_BLAKE:
case ALGO_BMW: case ALGO_BMW:
case ALGO_WHIRLPOOLX:
minmax = 0x40000000U; minmax = 0x40000000U;
break; break;
case ALGO_KECCAK:
case ALGO_LUFFA: case ALGO_LUFFA:
minmax = 0x2000000; case ALGO_SKEIN:
case ALGO_SKEIN2:
minmax = 0x1000000;
break; break;
case ALGO_C11: case ALGO_C11:
case ALGO_DEEP: case ALGO_DEEP:
case ALGO_HEAVY:
case ALGO_LYRA2v2: case ALGO_LYRA2v2:
case ALGO_S3: case ALGO_S3:
case ALGO_X11: case ALGO_X11:
@ -1729,7 +1735,6 @@ static void *miner_thread(void *userdata)
case ALGO_WHIRLPOOL: case ALGO_WHIRLPOOL:
minmax = 0x400000; minmax = 0x400000;
break; break;
case ALGO_KECCAK:
case ALGO_JACKPOT: case ALGO_JACKPOT:
case ALGO_X14: case ALGO_X14:
case ALGO_X15: case ALGO_X15:

4
cuda.cpp

@ -176,9 +176,11 @@ void cuda_reset_device(int thr_id, bool *init)
} }
cudaDeviceReset(); cudaDeviceReset();
if (opt_cudaschedule >= 0) { if (opt_cudaschedule >= 0) {
cudaSetDevice(dev_id);
cudaSetDeviceFlags((unsigned)(opt_cudaschedule & cudaDeviceScheduleMask)); cudaSetDeviceFlags((unsigned)(opt_cudaschedule & cudaDeviceScheduleMask));
} else {
cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
} }
cudaDeviceSynchronize();
} }
// return free memory in megabytes // return free memory in megabytes

11
cuda_checkhash.cu

@ -121,6 +121,9 @@ uint32_t cuda_check_hash(int thr_id, uint32_t threads, uint32_t startNounce, uin
dim3 grid((threads + threadsperblock - 1) / threadsperblock); dim3 grid((threads + threadsperblock - 1) / threadsperblock);
dim3 block(threadsperblock); dim3 block(threadsperblock);
if (bench_algo >= 0) // dont interrupt the global benchmark
return UINT32_MAX;
if (!init_done) { if (!init_done) {
applog(LOG_ERR, "missing call to cuda_check_cpu_init"); applog(LOG_ERR, "missing call to cuda_check_cpu_init");
return UINT32_MAX; return UINT32_MAX;
@ -143,6 +146,9 @@ uint32_t cuda_check_hash_32(int thr_id, uint32_t threads, uint32_t startNounce,
dim3 grid((threads + threadsperblock - 1) / threadsperblock); dim3 grid((threads + threadsperblock - 1) / threadsperblock);
dim3 block(threadsperblock); dim3 block(threadsperblock);
if (bench_algo >= 0) // dont interrupt the global benchmark
return UINT32_MAX;
if (!init_done) { if (!init_done) {
applog(LOG_ERR, "missing call to cuda_check_cpu_init"); applog(LOG_ERR, "missing call to cuda_check_cpu_init");
return UINT32_MAX; return UINT32_MAX;
@ -237,9 +243,12 @@ uint32_t cuda_check_hash_branch(int thr_id, uint32_t threads, uint32_t startNoun
uint32_t result = UINT32_MAX; uint32_t result = UINT32_MAX;
if (bench_algo >= 0) // dont interrupt the global benchmark
return result;
if (!init_done) { if (!init_done) {
applog(LOG_ERR, "missing call to cuda_check_cpu_init"); applog(LOG_ERR, "missing call to cuda_check_cpu_init");
return UINT32_MAX; return result;
} }
cudaMemset(d_resNonces[thr_id], 0xff, sizeof(uint32_t)); cudaMemset(d_resNonces[thr_id], 0xff, sizeof(uint32_t));

8
cuda_nist5.cu

@ -135,11 +135,17 @@ extern "C" int scanhash_nist5(int thr_id, struct work *work, uint32_t max_nonce,
} }
} }
if ((uint64_t) throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
break;
}
pdata[19] += throughput; pdata[19] += throughput;
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
out: out:
*hashes_done = pdata[19] - first_nonce;
#ifdef USE_STREAMS #ifdef USE_STREAMS
for (int i = 0; i < 5; i++) for (int i = 0; i < 5; i++)
cudaStreamDestroy(stream[i]); cudaStreamDestroy(stream[i]);

9
fuguecoin.cpp

@ -65,7 +65,9 @@ int scanhash_fugue256(int thr_id, struct work* work, uint32_t max_nonce, unsigne
uint32_t foundNounce = UINT32_MAX; uint32_t foundNounce = UINT32_MAX;
fugue256_cpu_hash(thr_id, throughput, pdata[19], NULL, &foundNounce); fugue256_cpu_hash(thr_id, throughput, pdata[19], NULL, &foundNounce);
if (foundNounce < UINT32_MAX) *hashes_done = pdata[19] - start_nonce + throughput;
if (foundNounce < UINT32_MAX && bench_algo < 0)
{ {
uint32_t vhash[8]; uint32_t vhash[8];
sph_fugue256_context ctx_fugue; sph_fugue256_context ctx_fugue;
@ -79,14 +81,13 @@ int scanhash_fugue256(int thr_id, struct work* work, uint32_t max_nonce, unsigne
{ {
work_set_target_ratio(work, vhash); work_set_target_ratio(work, vhash);
pdata[19] = foundNounce; pdata[19] = foundNounce;
*hashes_done = foundNounce - start_nonce + 1;
return 1; return 1;
} else { } else {
gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNounce); gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNounce);
} }
} }
if ((uint64_t) pdata[19] + throughput > (uint64_t) max_nonce) { if ((uint64_t) throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce; pdata[19] = max_nonce;
break; break;
} }
@ -95,7 +96,7 @@ int scanhash_fugue256(int thr_id, struct work* work, uint32_t max_nonce, unsigne
} while (!work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - start_nonce + 1; *hashes_done = pdata[19] - start_nonce;
return 0; return 0;
} }

11
groestlcoin.cpp

@ -39,7 +39,7 @@ int scanhash_groestlcoin(int thr_id, struct work *work, uint32_t max_nonce, unsi
uint32_t *outputHash = (uint32_t*)malloc((size_t) 64* throughput); uint32_t *outputHash = (uint32_t*)malloc((size_t) 64* throughput);
if (opt_benchmark) if (opt_benchmark)
ptarget[7] = 0x000ff; ptarget[7] = 0x001f;
if (!init[thr_id]) if (!init[thr_id])
{ {
@ -62,7 +62,7 @@ int scanhash_groestlcoin(int thr_id, struct work *work, uint32_t max_nonce, unsi
// GPU hash // GPU hash
groestlcoin_cpu_hash(thr_id, throughput, pdata[19], outputHash, &foundNounce); groestlcoin_cpu_hash(thr_id, throughput, pdata[19], outputHash, &foundNounce);
if (foundNounce < UINT32_MAX) if (foundNounce < UINT32_MAX && bench_algo < 0)
{ {
uint32_t _ALIGN(64) vhash[8]; uint32_t _ALIGN(64) vhash[8];
endiandata[19] = swab32(foundNounce); endiandata[19] = swab32(foundNounce);
@ -78,14 +78,15 @@ int scanhash_groestlcoin(int thr_id, struct work *work, uint32_t max_nonce, unsi
} }
} }
if ((uint64_t) pdata[19] + throughput > max_nonce) { if ((uint64_t) throughput + pdata[19] >= max_nonce) {
*hashes_done = pdata[19] - start_nonce + 1;
pdata[19] = max_nonce; pdata[19] = max_nonce;
break; break;
} }
pdata[19] += throughput; pdata[19] += throughput;
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - start_nonce;
free(outputHash); free(outputHash);
return 0; return 0;

575
heavy/heavy.cu

@ -43,99 +43,99 @@ extern uint32_t *heavy_heftyHashes[MAX_GPUS];
/* Combines top 64-bits from each hash into a single hash */ /* Combines top 64-bits from each hash into a single hash */
static void combine_hashes(uint32_t *out, const uint32_t *hash1, const uint32_t *hash2, const uint32_t *hash3, const uint32_t *hash4) static void combine_hashes(uint32_t *out, const uint32_t *hash1, const uint32_t *hash2, const uint32_t *hash3, const uint32_t *hash4)
{ {
const uint32_t *hash[4] = { hash1, hash2, hash3, hash4 }; const uint32_t *hash[4] = { hash1, hash2, hash3, hash4 };
int bits; int bits;
unsigned int i; unsigned int i;
uint32_t mask; uint32_t mask;
unsigned int k; unsigned int k;
/* Transpose first 64 bits of each hash into out */ /* Transpose first 64 bits of each hash into out */
memset(out, 0, 32); memset(out, 0, 32);
bits = 0; bits = 0;
for (i = 7; i >= 6; i--) { for (i = 7; i >= 6; i--) {
for (mask = 0x80000000; mask; mask >>= 1) { for (mask = 0x80000000; mask; mask >>= 1) {
for (k = 0; k < 4; k++) { for (k = 0; k < 4; k++) {
out[(255 - bits)/32] <<= 1; out[(255 - bits)/32] <<= 1;
if ((hash[k][i] & mask) != 0) if ((hash[k][i] & mask) != 0)
out[(255 - bits)/32] |= 1; out[(255 - bits)/32] |= 1;
bits++; bits++;
} }
} }
} }
} }
#ifdef _MSC_VER #ifdef _MSC_VER
#include <intrin.h> #include <intrin.h>
static uint32_t __inline bitsset( uint32_t x ) static uint32_t __inline bitsset( uint32_t x )
{ {
DWORD r = 0; DWORD r = 0;
_BitScanReverse(&r, x); _BitScanReverse(&r, x);
return r; return r;
} }
#else #else
static uint32_t bitsset( uint32_t x ) static uint32_t bitsset( uint32_t x )
{ {
return 31-__builtin_clz(x); return 31-__builtin_clz(x);
} }
#endif #endif
// Finde das high bit in einem Multiword-Integer. // Finde das high bit in einem Multiword-Integer.
static int findhighbit(const uint32_t *ptarget, int words) static int findhighbit(const uint32_t *ptarget, int words)
{ {
int i; int i;
int highbit = 0; int highbit = 0;
for (i=words-1; i >= 0; --i) for (i=words-1; i >= 0; --i)
{ {
if (ptarget[i] != 0) { if (ptarget[i] != 0) {
highbit = i*32 + bitsset(ptarget[i])+1; highbit = i*32 + bitsset(ptarget[i])+1;
break; break;
} }
} }
return highbit; return highbit;
} }
// Generiere ein Multiword-Integer das die Zahl // Generiere ein Multiword-Integer das die Zahl
// (2 << highbit) - 1 repräsentiert. // (2 << highbit) - 1 repräsentiert.
static void genmask(uint32_t *ptarget, int words, int highbit) static void genmask(uint32_t *ptarget, int words, int highbit)
{ {
int i; int i;
for (i=words-1; i >= 0; --i) for (i=words-1; i >= 0; --i)
{ {
if ((i+1)*32 <= highbit) if ((i+1)*32 <= highbit)
ptarget[i] = UINT32_MAX; ptarget[i] = UINT32_MAX;
else if (i*32 > highbit) else if (i*32 > highbit)
ptarget[i] = 0x00000000; ptarget[i] = 0x00000000;
else else
ptarget[i] = (1 << (highbit-i*32)) - 1; ptarget[i] = (1 << (highbit-i*32)) - 1;
} }
} }
struct check_nonce_for_remove struct check_nonce_for_remove
{ {
check_nonce_for_remove(uint64_t target, uint32_t *hashes, uint32_t hashlen, uint32_t startNonce) : check_nonce_for_remove(uint64_t target, uint32_t *hashes, uint32_t hashlen, uint32_t startNonce) :
m_target(target), m_target(target),
m_hashes(hashes), m_hashes(hashes),
m_hashlen(hashlen), m_hashlen(hashlen),
m_startNonce(startNonce) { } m_startNonce(startNonce) { }
uint64_t m_target; uint64_t m_target;
uint32_t *m_hashes; uint32_t *m_hashes;
uint32_t m_hashlen; uint32_t m_hashlen;
uint32_t m_startNonce; uint32_t m_startNonce;
__device__ __device__
bool operator()(const uint32_t x) bool operator()(const uint32_t x)
{ {
// Position im Hash Buffer // Position im Hash Buffer
uint32_t hashIndex = x - m_startNonce; uint32_t hashIndex = x - m_startNonce;
// Wert des Hashes (als uint64_t) auslesen. // Wert des Hashes (als uint64_t) auslesen.
// Steht im 6. und 7. Wort des Hashes (jeder dieser Hashes hat 512 Bits) // Steht im 6. und 7. Wort des Hashes (jeder dieser Hashes hat 512 Bits)
uint64_t hashValue = *((uint64_t*)(&m_hashes[m_hashlen*hashIndex + 6])); uint64_t hashValue = *((uint64_t*)(&m_hashes[m_hashlen*hashIndex + 6]));
bool res = (hashValue & m_target) != hashValue; bool res = (hashValue & m_target) != hashValue;
//printf("ndx=%x val=%08x target=%lx\n", hashIndex, hashValue, m_target); //printf("ndx=%x val=%08x target=%lx\n", hashIndex, hashValue, m_target);
// gegen das Target prüfen. Es dürfen nur Bits aus dem Target gesetzt sein. // gegen das Target prüfen. Es dürfen nur Bits aus dem Target gesetzt sein.
return res; return res;
} }
}; };
static bool init[MAX_GPUS] = { 0 }; static bool init[MAX_GPUS] = { 0 };
@ -143,245 +143,252 @@ static bool init[MAX_GPUS] = { 0 };
__host__ __host__
int scanhash_heavy(int thr_id, struct work *work, uint32_t max_nonce, unsigned long *hashes_done, uint32_t maxvote, int blocklen) int scanhash_heavy(int thr_id, struct work *work, uint32_t max_nonce, unsigned long *hashes_done, uint32_t maxvote, int blocklen)
{ {
uint32_t *pdata = work->data; uint32_t *pdata = work->data;
uint32_t *ptarget = work->target; uint32_t *ptarget = work->target;
const uint32_t first_nonce = pdata[19]; const uint32_t first_nonce = pdata[19];
// CUDA will process thousands of threads. // CUDA will process thousands of threads.
uint32_t throughput = cuda_default_throughput(thr_id, (1U << 19) - 256); uint32_t throughput = cuda_default_throughput(thr_id, (1U << 19) - 256);
if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce); if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);
int rc = 0; int rc = 0;
uint32_t *hash = NULL; uint32_t *hash = NULL;
uint32_t *cpu_nonceVector = NULL; uint32_t *cpu_nonceVector = NULL;
CUDA_SAFE_CALL(cudaMallocHost(&hash, throughput*8*sizeof(uint32_t)));
CUDA_SAFE_CALL(cudaMallocHost(&cpu_nonceVector, throughput*sizeof(uint32_t))); int nrmCalls[6];
memset(nrmCalls, 0, sizeof(int) * 6);
int nrmCalls[6];
memset(nrmCalls, 0, sizeof(int) * 6); if (opt_benchmark)
ptarget[7] = 0x000f;
if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x00ff; // für jeden Hash ein individuelles Target erstellen basierend
// auf dem höchsten Bit, das in ptarget gesetzt ist.
// für jeden Hash ein individuelles Target erstellen basierend int highbit = findhighbit(ptarget, 8);
// auf dem höchsten Bit, das in ptarget gesetzt ist. uint32_t target2[2], target3[2], target4[2], target5[2];
int highbit = findhighbit(ptarget, 8); genmask(target2, 2, highbit/4+(((highbit%4)>3)?1:0) ); // SHA256
uint32_t target2[2], target3[2], target4[2], target5[2]; genmask(target3, 2, highbit/4+(((highbit%4)>2)?1:0) ); // keccak512
genmask(target2, 2, highbit/4+(((highbit%4)>3)?1:0) ); // SHA256 genmask(target4, 2, highbit/4+(((highbit%4)>1)?1:0) ); // groestl512
genmask(target3, 2, highbit/4+(((highbit%4)>2)?1:0) ); // keccak512 genmask(target5, 2, highbit/4+(((highbit%4)>0)?1:0) ); // blake512
genmask(target4, 2, highbit/4+(((highbit%4)>1)?1:0) ); // groestl512
genmask(target5, 2, highbit/4+(((highbit%4)>0)?1:0) ); // blake512 if (!init[thr_id])
{
if (!init[thr_id]) cudaSetDevice(device_map[thr_id]);
{
hefty_cpu_init(thr_id, throughput); hefty_cpu_init(thr_id, throughput);
sha256_cpu_init(thr_id, throughput); sha256_cpu_init(thr_id, throughput);
keccak512_cpu_init(thr_id, throughput); keccak512_cpu_init(thr_id, throughput);
groestl512_cpu_init(thr_id, throughput); groestl512_cpu_init(thr_id, throughput);
blake512_cpu_init(thr_id, throughput); blake512_cpu_init(thr_id, throughput);
combine_cpu_init(thr_id, throughput); combine_cpu_init(thr_id, throughput);
CUDA_SAFE_CALL(cudaMalloc(&heavy_nonceVector[thr_id], sizeof(uint32_t) * throughput)); CUDA_SAFE_CALL(cudaMalloc(&heavy_nonceVector[thr_id], sizeof(uint32_t) * throughput));
init[thr_id] = true; init[thr_id] = true;
} }
if (blocklen == HEAVYCOIN_BLKHDR_SZ) // weird but require at least one cudaSetDevice first
{ CUDA_SAFE_CALL(cudaMallocHost(&hash, (size_t) 32 * throughput));
uint16_t *ext = (uint16_t *)&pdata[20]; CUDA_SAFE_CALL(cudaMallocHost(&cpu_nonceVector, sizeof(uint32_t) * throughput));
if (opt_vote > maxvote && !opt_benchmark) { if (blocklen == HEAVYCOIN_BLKHDR_SZ)
applog(LOG_WARNING, "Your block reward vote (%hu) exceeds " {
"the maxvote reported by the pool (%hu).", uint16_t *ext = (uint16_t*) &pdata[20];
opt_vote, maxvote);
} if (opt_vote > maxvote && !opt_benchmark) {
applog(LOG_WARNING, "Your block reward vote (%hu) exceeds the maxvote reported by the pool (%hu).",
if (opt_trust_pool && opt_vote > maxvote) { opt_vote, maxvote);
applog(LOG_WARNING, "Capping block reward vote to maxvote reported by pool."); }
ext[0] = maxvote;
} if (opt_trust_pool && opt_vote > maxvote) {
else applog(LOG_WARNING, "Capping block reward vote to maxvote reported by pool.");
ext[0] = opt_vote; ext[0] = maxvote;
} }
else
// Setze die Blockdaten ext[0] = opt_vote;
hefty_cpu_setBlock(thr_id, throughput, pdata, blocklen); }
sha256_cpu_setBlock(pdata, blocklen);
keccak512_cpu_setBlock(pdata, blocklen); // Setze die Blockdaten
groestl512_cpu_setBlock(pdata, blocklen); hefty_cpu_setBlock(thr_id, throughput, pdata, blocklen);
blake512_cpu_setBlock(pdata, blocklen); sha256_cpu_setBlock(pdata, blocklen);
keccak512_cpu_setBlock(pdata, blocklen);
do { groestl512_cpu_setBlock(pdata, blocklen);
uint32_t actualNumberOfValuesInNonceVectorGPU = throughput; blake512_cpu_setBlock(pdata, blocklen);
////// Compaction init do {
uint32_t actualNumberOfValuesInNonceVectorGPU = throughput;
hefty_cpu_hash(thr_id, throughput, pdata[19]);
sha256_cpu_hash(thr_id, throughput, pdata[19]); ////// Compaction init
// Hier ist die längste CPU Wartephase. Deshalb ein strategisches MyStreamSynchronize() hier. hefty_cpu_hash(thr_id, throughput, pdata[19]);
MyStreamSynchronize(NULL, 1, thr_id); sha256_cpu_hash(thr_id, throughput, pdata[19]);
// Hier ist die längste CPU Wartephase. Deshalb ein strategisches MyStreamSynchronize() hier.
MyStreamSynchronize(NULL, 1, thr_id);
#if USE_THRUST #if USE_THRUST
thrust::device_ptr<uint32_t> devNoncePtr(heavy_nonceVector[thr_id]); thrust::device_ptr<uint32_t> devNoncePtr(heavy_nonceVector[thr_id]);
thrust::device_ptr<uint32_t> devNoncePtrEnd((heavy_nonceVector[thr_id]) + throughput); thrust::device_ptr<uint32_t> devNoncePtrEnd((heavy_nonceVector[thr_id]) + throughput);
////// Compaction ////// Compaction
uint64_t *t = (uint64_t*) target2; uint64_t *t = (uint64_t*) target2;
devNoncePtrEnd = thrust::remove_if(devNoncePtr, devNoncePtrEnd, check_nonce_for_remove(*t, d_hash2output[thr_id], 8, pdata[19])); devNoncePtrEnd = thrust::remove_if(devNoncePtr, devNoncePtrEnd, check_nonce_for_remove(*t, d_hash2output[thr_id], 8, pdata[19]));
actualNumberOfValuesInNonceVectorGPU = (uint32_t)(devNoncePtrEnd - devNoncePtr); actualNumberOfValuesInNonceVectorGPU = (uint32_t)(devNoncePtrEnd - devNoncePtr);
if(actualNumberOfValuesInNonceVectorGPU == 0) if(actualNumberOfValuesInNonceVectorGPU == 0)
goto emptyNonceVector; goto emptyNonceVector;
keccak512_cpu_hash(thr_id, actualNumberOfValuesInNonceVectorGPU, pdata[19]); keccak512_cpu_hash(thr_id, actualNumberOfValuesInNonceVectorGPU, pdata[19]);
////// Compaction ////// Compaction
t = (uint64_t*) target3; t = (uint64_t*) target3;
devNoncePtrEnd = thrust::remove_if(devNoncePtr, devNoncePtrEnd, check_nonce_for_remove(*t, d_hash3output[thr_id], 16, pdata[19])); devNoncePtrEnd = thrust::remove_if(devNoncePtr, devNoncePtrEnd, check_nonce_for_remove(*t, d_hash3output[thr_id], 16, pdata[19]));
actualNumberOfValuesInNonceVectorGPU = (uint32_t)(devNoncePtrEnd - devNoncePtr); actualNumberOfValuesInNonceVectorGPU = (uint32_t)(devNoncePtrEnd - devNoncePtr);
if(actualNumberOfValuesInNonceVectorGPU == 0) if(actualNumberOfValuesInNonceVectorGPU == 0)
goto emptyNonceVector; goto emptyNonceVector;
blake512_cpu_hash(thr_id, actualNumberOfValuesInNonceVectorGPU, pdata[19]); blake512_cpu_hash(thr_id, actualNumberOfValuesInNonceVectorGPU, pdata[19]);
////// Compaction ////// Compaction
t = (uint64_t*) target5; t = (uint64_t*) target5;
devNoncePtrEnd = thrust::remove_if(devNoncePtr, devNoncePtrEnd, check_nonce_for_remove(*t, d_hash5output[thr_id], 16, pdata[19])); devNoncePtrEnd = thrust::remove_if(devNoncePtr, devNoncePtrEnd, check_nonce_for_remove(*t, d_hash5output[thr_id], 16, pdata[19]));
actualNumberOfValuesInNonceVectorGPU = (uint32_t)(devNoncePtrEnd - devNoncePtr); actualNumberOfValuesInNonceVectorGPU = (uint32_t)(devNoncePtrEnd - devNoncePtr);
if(actualNumberOfValuesInNonceVectorGPU == 0) if(actualNumberOfValuesInNonceVectorGPU == 0)
goto emptyNonceVector; goto emptyNonceVector;
groestl512_cpu_hash(thr_id, actualNumberOfValuesInNonceVectorGPU, pdata[19]); groestl512_cpu_hash(thr_id, actualNumberOfValuesInNonceVectorGPU, pdata[19]);
////// Compaction ////// Compaction
t = (uint64_t*) target4; t = (uint64_t*) target4;
devNoncePtrEnd = thrust::remove_if(devNoncePtr, devNoncePtrEnd, check_nonce_for_remove(*t, d_hash4output[thr_id], 16, pdata[19])); devNoncePtrEnd = thrust::remove_if(devNoncePtr, devNoncePtrEnd, check_nonce_for_remove(*t, d_hash4output[thr_id], 16, pdata[19]));
actualNumberOfValuesInNonceVectorGPU = (uint32_t)(devNoncePtrEnd - devNoncePtr); actualNumberOfValuesInNonceVectorGPU = (uint32_t)(devNoncePtrEnd - devNoncePtr);
#else #else
// todo // todo (nvlabs cub ?)
actualNumberOfValuesInNonceVectorGPU = 0; actualNumberOfValuesInNonceVectorGPU = 0;
#endif #endif
if(actualNumberOfValuesInNonceVectorGPU == 0) if(actualNumberOfValuesInNonceVectorGPU == 0)
goto emptyNonceVector; goto emptyNonceVector;
// combine // combine
combine_cpu_hash(thr_id, actualNumberOfValuesInNonceVectorGPU, pdata[19], hash); combine_cpu_hash(thr_id, actualNumberOfValuesInNonceVectorGPU, pdata[19], hash);
if (opt_tracegpu) { if (opt_tracegpu) {
applog(LOG_BLUE, "heavy GPU hash:"); applog(LOG_BLUE, "heavy GPU hash:");
applog_hash((uchar*)hash); applog_hash((uchar*)hash);
} }
// Ergebnisse kopieren // Ergebnisse kopieren
if(actualNumberOfValuesInNonceVectorGPU > 0) if(actualNumberOfValuesInNonceVectorGPU > 0)
{ {
size_t size = sizeof(uint32_t) * actualNumberOfValuesInNonceVectorGPU; size_t size = sizeof(uint32_t) * actualNumberOfValuesInNonceVectorGPU;
CUDA_SAFE_CALL(cudaMemcpy(cpu_nonceVector, heavy_nonceVector[thr_id], size, cudaMemcpyDeviceToHost)); cudaMemcpy(cpu_nonceVector, heavy_nonceVector[thr_id], size, cudaMemcpyDeviceToHost);
cudaThreadSynchronize();
for (uint32_t i=0; i < actualNumberOfValuesInNonceVectorGPU; i++)
for (uint32_t i=0; i < actualNumberOfValuesInNonceVectorGPU; i++) {
{ uint32_t nonce = cpu_nonceVector[i];
uint32_t nonce = cpu_nonceVector[i]; uint32_t *foundhash = &hash[8*i];
uint32_t *foundhash = &hash[8*i]; if (foundhash[7] <= ptarget[7] && fulltest(foundhash, ptarget)) {
if (foundhash[7] <= ptarget[7] && fulltest(foundhash, ptarget)) { uint32_t vhash[8];
uint32_t vhash[8]; pdata[19] += nonce - pdata[19];
pdata[19] += nonce - pdata[19]; heavycoin_hash((uchar*)vhash, (uchar*)pdata, blocklen);
heavycoin_hash((uchar*)vhash, (uchar*)pdata, blocklen); if (memcmp(vhash, foundhash, 32)) {
if (memcmp(vhash, foundhash, 32)) { gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", nonce);
gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", nonce); } else {
} else { work_set_target_ratio(work, vhash);
*hashes_done = pdata[19] - first_nonce; rc = 1;
work_set_target_ratio(work, vhash); goto exit;
rc = 1; }
goto exit; }
} }
} }
}
}
emptyNonceVector: emptyNonceVector:
if ((uint64_t) throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
break;
}
pdata[19] += throughput;
pdata[19] += throughput; } while (!work_restart[thr_id].restart);
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce;
exit: exit:
cudaFreeHost(cpu_nonceVector); *hashes_done = pdata[19] - first_nonce;
cudaFreeHost(hash);
return rc; cudaFreeHost(cpu_nonceVector);
cudaFreeHost(hash);
CUDA_LOG_ERROR();
return rc;
} }
// cleanup // cleanup
extern "C" void free_heavy(int thr_id) extern "C" void free_heavy(int thr_id)
{ {
if (!init[thr_id]) if (!init[thr_id])
return; return;
cudaThreadSynchronize(); cudaThreadSynchronize();
cudaFree(heavy_nonceVector[thr_id]); cudaFree(heavy_nonceVector[thr_id]);
blake512_cpu_free(thr_id); blake512_cpu_free(thr_id);
groestl512_cpu_free(thr_id); groestl512_cpu_free(thr_id);
hefty_cpu_free(thr_id); hefty_cpu_free(thr_id);
keccak512_cpu_free(thr_id); keccak512_cpu_free(thr_id);
sha256_cpu_free(thr_id); sha256_cpu_free(thr_id);
combine_cpu_free(thr_id); combine_cpu_free(thr_id);
init[thr_id] = false; init[thr_id] = false;
cudaDeviceSynchronize(); cudaDeviceSynchronize();
} }
__host__ __host__
void heavycoin_hash(uchar* output, const uchar* input, int len) void heavycoin_hash(uchar* output, const uchar* input, int len)
{ {
unsigned char hash1[32]; unsigned char hash1[32];
unsigned char hash2[32]; unsigned char hash2[32];
uint32_t hash3[16]; uint32_t hash3[16];
uint32_t hash4[16]; uint32_t hash4[16];
uint32_t hash5[16]; uint32_t hash5[16];
uint32_t *final; uint32_t *final;
SHA256_CTX ctx; SHA256_CTX ctx;
sph_keccak512_context keccakCtx; sph_keccak512_context keccakCtx;
sph_groestl512_context groestlCtx; sph_groestl512_context groestlCtx;
sph_blake512_context blakeCtx; sph_blake512_context blakeCtx;
HEFTY1(input, len, hash1); HEFTY1(input, len, hash1);
/* HEFTY1 is new, so take an extra security measure to eliminate /* HEFTY1 is new, so take an extra security measure to eliminate
* the possiblity of collisions: * the possiblity of collisions:
* *
* Hash(x) = SHA256(x + HEFTY1(x)) * Hash(x) = SHA256(x + HEFTY1(x))
* *
* N.B. '+' is concatenation. * N.B. '+' is concatenation.
*/ */
SHA256_Init(&ctx); SHA256_Init(&ctx);
SHA256_Update(&ctx, input, len); SHA256_Update(&ctx, input, len);
SHA256_Update(&ctx, hash1, sizeof(hash1)); SHA256_Update(&ctx, hash1, sizeof(hash1));
SHA256_Final(hash2, &ctx); SHA256_Final(hash2, &ctx);
/* Additional security: Do not rely on a single cryptographic hash /* Additional security: Do not rely on a single cryptographic hash
* function. Instead, combine the outputs of 4 of the most secure * function. Instead, combine the outputs of 4 of the most secure
* cryptographic hash functions-- SHA256, KECCAK512, GROESTL512 * cryptographic hash functions-- SHA256, KECCAK512, GROESTL512
* and BLAKE512. * and BLAKE512.
*/ */
sph_keccak512_init(&keccakCtx); sph_keccak512_init(&keccakCtx);
sph_keccak512(&keccakCtx, input, len); sph_keccak512(&keccakCtx, input, len);
sph_keccak512(&keccakCtx, hash1, sizeof(hash1)); sph_keccak512(&keccakCtx, hash1, sizeof(hash1));
sph_keccak512_close(&keccakCtx, (void *)&hash3); sph_keccak512_close(&keccakCtx, (void *)&hash3);
sph_groestl512_init(&groestlCtx); sph_groestl512_init(&groestlCtx);
sph_groestl512(&groestlCtx, input, len); sph_groestl512(&groestlCtx, input, len);
sph_groestl512(&groestlCtx, hash1, sizeof(hash1)); sph_groestl512(&groestlCtx, hash1, sizeof(hash1));
sph_groestl512_close(&groestlCtx, (void *)&hash4); sph_groestl512_close(&groestlCtx, (void *)&hash4);
sph_blake512_init(&blakeCtx); sph_blake512_init(&blakeCtx);
sph_blake512(&blakeCtx, input, len); sph_blake512(&blakeCtx, input, len);
sph_blake512(&blakeCtx, (unsigned char *)&hash1, sizeof(hash1)); sph_blake512(&blakeCtx, (unsigned char *)&hash1, sizeof(hash1));
sph_blake512_close(&blakeCtx, (void *)&hash5); sph_blake512_close(&blakeCtx, (void *)&hash5);
final = (uint32_t *)output; final = (uint32_t *)output;
combine_hashes(final, (uint32_t *)hash2, hash3, hash4, hash5); combine_hashes(final, (uint32_t *)hash2, hash3, hash4, hash5);
} }

11
lyra2/lyra2RE.cu

@ -122,14 +122,14 @@ extern "C" int scanhash_lyra2(int thr_id, struct work* work, uint32_t max_nonce,
int order = 0; int order = 0;
uint32_t foundNonce; uint32_t foundNonce;
*hashes_done = pdata[19] - first_nonce + throughput;
blake256_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++); blake256_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
keccak256_cpu_hash_32(thr_id, throughput, pdata[19], d_hash[thr_id], order++); keccak256_cpu_hash_32(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
lyra2_cpu_hash_32(thr_id, throughput, pdata[19], d_hash[thr_id], order++); lyra2_cpu_hash_32(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
skein256_cpu_hash_32(thr_id, throughput, pdata[19], d_hash[thr_id], order++); skein256_cpu_hash_32(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
TRACE("S") TRACE("S")
*hashes_done = pdata[19] - first_nonce + throughput;
foundNonce = groestl256_cpu_hash_32(thr_id, throughput, pdata[19], d_hash[thr_id], order++); foundNonce = groestl256_cpu_hash_32(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
if (foundNonce != UINT32_MAX) if (foundNonce != UINT32_MAX)
{ {
@ -162,10 +162,15 @@ extern "C" int scanhash_lyra2(int thr_id, struct work* work, uint32_t max_nonce,
} }
} }
if ((uint64_t)throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
break;
}
pdata[19] += throughput; pdata[19] += throughput;
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce;
return 0; return 0;
} }

11
lyra2/lyra2REv2.cu

@ -96,14 +96,13 @@ extern "C" int scanhash_lyra2v2(int thr_id, struct work* work, uint32_t max_nonc
cudaDeviceReset(); cudaDeviceReset();
// reduce cpu usage // reduce cpu usage
cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
CUDA_LOG_ERROR();
} }
CUDA_LOG_ERROR();
blake256_cpu_init(thr_id, throughput); blake256_cpu_init(thr_id, throughput);
keccak256_cpu_init(thr_id,throughput); keccak256_cpu_init(thr_id,throughput);
skein256_cpu_init(thr_id, throughput); skein256_cpu_init(thr_id, throughput);
bmw256_cpu_init(thr_id, throughput); bmw256_cpu_init(thr_id, throughput);
CUDA_LOG_ERROR();
// SM 3 implentation requires a bit more memory // SM 3 implentation requires a bit more memory
if (device_sm[dev_id] < 500 || cuda_arch[dev_id] < 500) if (device_sm[dev_id] < 500 || cuda_arch[dev_id] < 500)
@ -175,11 +174,15 @@ extern "C" int scanhash_lyra2v2(int thr_id, struct work* work, uint32_t max_nonc
} }
} }
if ((uint64_t)throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
break;
}
pdata[19] += throughput; pdata[19] += throughput;
} while (!work_restart[thr_id].restart && (max_nonce > ((uint64_t)(pdata[19]) + throughput))); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce + 1; *hashes_done = pdata[19] - first_nonce;
return 0; return 0;
} }

13
myriadgroestl.cpp

@ -43,7 +43,7 @@ int scanhash_myriad(int thr_id, struct work *work, uint32_t max_nonce, unsigned
uint32_t *outputHash = (uint32_t*)malloc(throughput * 64); uint32_t *outputHash = (uint32_t*)malloc(throughput * 64);
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x0000ff; ptarget[7] = 0x0000ff;
// init // init
if(!init[thr_id]) if(!init[thr_id])
@ -63,11 +63,11 @@ int scanhash_myriad(int thr_id, struct work *work, uint32_t max_nonce, unsigned
// GPU // GPU
uint32_t foundNounce = UINT32_MAX; uint32_t foundNounce = UINT32_MAX;
*hashes_done = pdata[19] - start_nonce + throughput;
myriadgroestl_cpu_hash(thr_id, throughput, pdata[19], outputHash, &foundNounce); myriadgroestl_cpu_hash(thr_id, throughput, pdata[19], outputHash, &foundNounce);
if (foundNounce < UINT32_MAX) *hashes_done = pdata[19] - start_nonce + throughput;
if (foundNounce < UINT32_MAX && bench_algo < 0)
{ {
uint32_t _ALIGN(64) vhash[8]; uint32_t _ALIGN(64) vhash[8];
endiandata[19] = swab32(foundNounce); endiandata[19] = swab32(foundNounce);
@ -82,8 +82,7 @@ int scanhash_myriad(int thr_id, struct work *work, uint32_t max_nonce, unsigned
} }
} }
if ((uint64_t) pdata[19] + throughput > max_nonce) { if ((uint64_t) throughput + pdata[19] >= max_nonce) {
*hashes_done = pdata[19] - start_nonce;
pdata[19] = max_nonce; pdata[19] = max_nonce;
break; break;
} }
@ -91,6 +90,8 @@ int scanhash_myriad(int thr_id, struct work *work, uint32_t max_nonce, unsigned
} while (!work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*hashes_done = max_nonce - start_nonce;
free(outputHash); free(outputHash);
return 0; return 0;
} }

9
neoscrypt/neoscrypt.cpp

@ -79,11 +79,16 @@ int scanhash_neoscrypt(int thr_id, struct work* work, uint32_t max_nonce, unsign
} }
} }
if ((uint64_t)throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
break;
}
pdata[19] += throughput; pdata[19] += throughput;
} while (!work_restart[thr_id].restart && (max_nonce > ((uint64_t)(pdata[19]) + throughput))); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce + 1; *hashes_done = pdata[19] - first_nonce;
return 0; return 0;
} }

7
pentablake.cu

@ -110,9 +110,14 @@ extern "C" int scanhash_pentablake(int thr_id, struct work *work, uint32_t max_n
} }
} }
if ((uint64_t) throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
break;
}
pdata[19] += throughput; pdata[19] += throughput;
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
return rc; return rc;
} }

7
qubit/deep.cu

@ -117,9 +117,14 @@ extern "C" int scanhash_deep(int thr_id, struct work* work, uint32_t max_nonce,
} }
} }
if ((uint64_t)throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
break;
}
pdata[19] += throughput; pdata[19] += throughput;
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce + 1; *hashes_done = pdata[19] - first_nonce + 1;
return 0; return 0;

25
qubit/luffa.cu

@ -36,19 +36,21 @@ extern "C" int scanhash_luffa(int thr_id, struct work* work, uint32_t max_nonce,
uint32_t *pdata = work->data; uint32_t *pdata = work->data;
uint32_t *ptarget = work->target; uint32_t *ptarget = work->target;
const uint32_t first_nonce = pdata[19]; const uint32_t first_nonce = pdata[19];
uint32_t throughput = cuda_default_throughput(thr_id, 1U << 22); // 256*256*8*8 uint32_t throughput = cuda_default_throughput(thr_id, 1U << 21);
if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce); if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x0000f; ptarget[7] = 0x0000f;
if (!init[thr_id]) if (!init[thr_id])
{ {
cudaSetDevice(device_map[thr_id]); cudaSetDevice(device_map[thr_id]);
CUDA_LOG_ERROR(); if (opt_cudaschedule == -1 && gpu_threads == 1) {
//if (opt_cudaschedule == -1) // to reduce cpu usage... cudaDeviceReset();
// cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); // reduce cpu usage
//CUDA_LOG_ERROR(); cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
CUDA_LOG_ERROR();
}
CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], (size_t) 64 * throughput)); CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], (size_t) 64 * throughput));
@ -65,10 +67,9 @@ extern "C" int scanhash_luffa(int thr_id, struct work* work, uint32_t max_nonce,
cuda_check_cpu_setTarget(ptarget); cuda_check_cpu_setTarget(ptarget);
do { do {
int order = 0; qubit_luffa512_cpu_hash_80(thr_id, (int) throughput, pdata[19], d_hash[thr_id], 0);
*hashes_done = pdata[19] - first_nonce + throughput;
qubit_luffa512_cpu_hash_80(thr_id, (int) throughput, pdata[19], 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]); uint32_t foundNonce = cuda_check_hash(thr_id, throughput, pdata[19], d_hash[thr_id]);
if (foundNonce != UINT32_MAX) if (foundNonce != UINT32_MAX)
@ -86,8 +87,8 @@ extern "C" int scanhash_luffa(int thr_id, struct work* work, uint32_t max_nonce,
} }
} }
if ((uint64_t) throughput + pdata[19] > max_nonce) { if ((uint64_t) throughput + pdata[19] >= max_nonce) {
// pdata[19] = max_nonce; pdata[19] = max_nonce;
break; break;
} }
@ -95,7 +96,7 @@ extern "C" int scanhash_luffa(int thr_id, struct work* work, uint32_t max_nonce,
} while (!work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce + 1; *hashes_done = pdata[19] - first_nonce;
return 0; return 0;
} }

6
skein.cu

@ -453,9 +453,7 @@ extern "C" int scanhash_skeincoin(int thr_id, struct work* work, uint32_t max_no
} }
} }
if ((uint64_t) throughput + pdata[19] > max_nonce) { if ((uint64_t) throughput + pdata[19] >= max_nonce) {
//applog(LOG_DEBUG, "done... max=%u", max_nonce);
*hashes_done = pdata[19] - first_nonce;
pdata[19] = max_nonce; pdata[19] = max_nonce;
break; break;
} }
@ -464,6 +462,8 @@ extern "C" int scanhash_skeincoin(int thr_id, struct work* work, uint32_t max_no
} while (!work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce;
return 0; return 0;
} }

5
skein2.cpp

@ -110,8 +110,7 @@ int scanhash_skein2(int thr_id, struct work* work, uint32_t max_nonce, unsigned
} }
} }
if ((uint64_t) throughput + pdata[19] > max_nonce) { if ((uint64_t) throughput + pdata[19] >= max_nonce) {
*hashes_done = pdata[19] - first_nonce;
pdata[19] = max_nonce; pdata[19] = max_nonce;
break; break;
} }
@ -120,6 +119,8 @@ int scanhash_skein2(int thr_id, struct work* work, uint32_t max_nonce, unsigned
} while (!work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce;
return 0; return 0;
} }

12
x11/c11.cu

@ -200,14 +200,20 @@ extern "C" int scanhash_c11(int thr_id, struct work* work, uint32_t max_nonce, u
} else { } else {
gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce); gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
pdata[19] = foundNonce + 1; pdata[19] = foundNonce + 1;
continue;
} }
} }
if ((uint64_t) throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
break;
}
pdata[19] += throughput; pdata[19] += throughput;
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce + 1; *hashes_done = pdata[19] - first_nonce;
return 0; return 0;
} }
@ -228,4 +234,4 @@ extern "C" void free_c11(int thr_id)
init[thr_id] = false; init[thr_id] = false;
cudaDeviceSynchronize(); cudaDeviceSynchronize();
} }

11
x11/fresh.cu

@ -78,18 +78,19 @@ extern "C" int scanhash_fresh(int thr_id, struct work* work, uint32_t max_nonce,
if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce); if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x00ff; ptarget[7] = 0x00ff;
if (!init[thr_id]) if (!init[thr_id])
{ {
cudaSetDevice(device_map[thr_id]); cudaSetDevice(device_map[thr_id]);
CUDA_LOG_ERROR();
CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], (size_t)64 * throughput + 4), -1);
x11_shavite512_cpu_init(thr_id, throughput); x11_shavite512_cpu_init(thr_id, throughput);
x11_simd512_cpu_init(thr_id, throughput); x11_simd512_cpu_init(thr_id, throughput);
x11_echo512_cpu_init(thr_id, throughput); x11_echo512_cpu_init(thr_id, throughput);
CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput + 4), 0);
cuda_check_cpu_init(thr_id, throughput); cuda_check_cpu_init(thr_id, throughput);
init[thr_id] = true; init[thr_id] = true;
@ -101,8 +102,6 @@ extern "C" int scanhash_fresh(int thr_id, struct work* work, uint32_t max_nonce,
x11_shavite512_setBlock_80((void*)endiandata); x11_shavite512_setBlock_80((void*)endiandata);
cuda_check_cpu_setTarget(ptarget); cuda_check_cpu_setTarget(ptarget);
do { do {
uint32_t Htarg = ptarget[7];
uint32_t foundNonce; uint32_t foundNonce;
int order = 0; int order = 0;
@ -128,7 +127,7 @@ extern "C" int scanhash_fresh(int thr_id, struct work* work, uint32_t max_nonce,
be32enc(&endiandata[19], foundNonce); be32enc(&endiandata[19], foundNonce);
fresh_hash(vhash64, endiandata); fresh_hash(vhash64, endiandata);
if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) { if (vhash64[7] <= ptarget[7] && fulltest(vhash64, ptarget)) {
int res = 1; int res = 1;
uint32_t secNonce = cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], 1); uint32_t secNonce = cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], 1);
work_set_target_ratio(work, vhash64); work_set_target_ratio(work, vhash64);

9
x11/s3.cu

@ -147,11 +147,16 @@ extern "C" int scanhash_s3(int thr_id, struct work* work, uint32_t max_nonce, un
} }
} }
if ((uint64_t) throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
break;
}
pdata[19] += throughput; pdata[19] += throughput;
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce + 1; *hashes_done = pdata[19] - first_nonce;
return 0; return 0;
} }

11
x11/x11.cu

@ -138,7 +138,7 @@ extern "C" int scanhash_x11(int thr_id, struct work* work, uint32_t max_nonce, u
if (x11_simd512_cpu_init(thr_id, throughput) != 0) { if (x11_simd512_cpu_init(thr_id, throughput) != 0) {
return 0; return 0;
} }
CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], 64 * throughput), 0); // why 64 ? CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], (size_t) 64 * throughput), 0);
cuda_check_cpu_init(thr_id, throughput); cuda_check_cpu_init(thr_id, throughput);
@ -205,14 +205,19 @@ extern "C" int scanhash_x11(int thr_id, struct work* work, uint32_t max_nonce, u
} else { } else {
gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce); gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce);
pdata[19] = foundNonce + 1; pdata[19] = foundNonce + 1;
continue;
} }
} }
if ((uint64_t) throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
break;
}
pdata[19] += throughput; pdata[19] += throughput;
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce + 1; *hashes_done = pdata[19] - first_nonce;
return 0; return 0;
} }

4
x13/x13.cu

@ -206,7 +206,7 @@ extern "C" int scanhash_x13(int thr_id, struct work* work, uint32_t max_nonce, u
} }
} }
if ((uint64_t)pdata[19] + throughput > (uint64_t)max_nonce) { if ((uint64_t)throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce; pdata[19] = max_nonce;
break; break;
} }
@ -214,7 +214,7 @@ extern "C" int scanhash_x13(int thr_id, struct work* work, uint32_t max_nonce, u
} while (!work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce + 1; *hashes_done = pdata[19] - first_nonce;
CUDA_LOG_ERROR(); CUDA_LOG_ERROR();

11
x15/whirlpool.cu

@ -83,12 +83,14 @@ extern "C" int scanhash_whirl(int thr_id, struct work* work, uint32_t max_nonce,
uint32_t foundNonce; uint32_t foundNonce;
int order = 0; int order = 0;
*hashes_done = pdata[19] - first_nonce + throughput;
whirlpool512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++); whirlpool512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
x15_whirlpool_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); x15_whirlpool_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
x15_whirlpool_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); x15_whirlpool_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
foundNonce = whirlpool512_cpu_finalhash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++); foundNonce = whirlpool512_cpu_finalhash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
if (foundNonce != UINT32_MAX) if (foundNonce != UINT32_MAX && bench_algo < 0)
{ {
const uint32_t Htarg = ptarget[7]; const uint32_t Htarg = ptarget[7];
uint32_t vhash[8]; uint32_t vhash[8];
@ -97,7 +99,6 @@ extern "C" int scanhash_whirl(int thr_id, struct work* work, uint32_t max_nonce,
if (vhash[7] <= Htarg && fulltest(vhash, ptarget)) { if (vhash[7] <= Htarg && fulltest(vhash, ptarget)) {
int res = 1; int res = 1;
*hashes_done = pdata[19] - first_nonce + throughput;
work_set_target_ratio(work, vhash); work_set_target_ratio(work, vhash);
#if 0 #if 0
uint32_t secNonce = cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], 1); uint32_t secNonce = cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], 1);
@ -112,9 +113,13 @@ extern "C" int scanhash_whirl(int thr_id, struct work* work, uint32_t max_nonce,
applog(LOG_WARNING, "GPU #%d: result for %08x does not validate on CPU!", device_map[thr_id], foundNonce); applog(LOG_WARNING, "GPU #%d: result for %08x does not validate on CPU!", device_map[thr_id], foundNonce);
} }
} }
if ((uint64_t) throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
break;
}
pdata[19] += throughput; pdata[19] += throughput;
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce; *hashes_done = pdata[19] - first_nonce;
return 0; return 0;

20
x15/whirlpoolx.cu

@ -49,12 +49,12 @@ extern "C" int scanhash_whirlx(int thr_id, struct work* work, uint32_t max_nonc
if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce); if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x0000ff; ptarget[7] = 0x000f;
if (!init[thr_id]) { if (!init[thr_id]) {
cudaSetDevice(device_map[thr_id]); cudaSetDevice(device_map[thr_id]);
CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], (size_t) 64 * throughput), 0); CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], (size_t) 64 * throughput), -1);
whirlpoolx_cpu_init(thr_id, throughput); whirlpoolx_cpu_init(thr_id, throughput);
@ -70,15 +70,16 @@ extern "C" int scanhash_whirlx(int thr_id, struct work* work, uint32_t max_nonc
do { do {
uint32_t foundNonce = whirlpoolx_cpu_hash(thr_id, throughput, pdata[19]); uint32_t foundNonce = whirlpoolx_cpu_hash(thr_id, throughput, pdata[19]);
if (foundNonce != UINT32_MAX)
*(hashes_done) = pdata[19] - first_nonce + throughput;
if (foundNonce != UINT32_MAX && bench_algo < 0)
{ {
const uint32_t Htarg = ptarget[7]; const uint32_t Htarg = ptarget[7];
uint32_t vhash64[8]; uint32_t vhash64[8];
be32enc(&endiandata[19], foundNonce); be32enc(&endiandata[19], foundNonce);
whirlxHash(vhash64, endiandata); whirlxHash(vhash64, endiandata);
*hashes_done = pdata[19] - first_nonce + throughput;
if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) { if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) {
work_set_target_ratio(work, vhash64); work_set_target_ratio(work, vhash64);
pdata[19] = foundNonce; pdata[19] = foundNonce;
@ -88,15 +89,16 @@ extern "C" int scanhash_whirlx(int thr_id, struct work* work, uint32_t max_nonc
} }
} }
pdata[19] += throughput; if ((uint64_t)throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
if (((uint64_t)pdata[19]+throughput) >= max_nonce) {
break; break;
} }
pdata[19] += throughput;
} while (!work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*(hashes_done) = pdata[19] - first_nonce + 1; *(hashes_done) = pdata[19] - first_nonce;
return 0; return 0;
} }

10
x15/x14.cu

@ -219,13 +219,19 @@ extern "C" int scanhash_x14(int thr_id, struct work* work, uint32_t max_nonce,
gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", foundNonce); 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; pdata[19] += throughput;
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
CUDA_LOG_ERROR(); CUDA_LOG_ERROR();
*hashes_done = pdata[19] - first_nonce + 1; *hashes_done = pdata[19] - first_nonce;
return 0; return 0;
} }

9
x15/x15.cu

@ -226,11 +226,16 @@ extern "C" int scanhash_x15(int thr_id, struct work* work, uint32_t max_nonce,
} }
} }
if ((uint64_t) throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
break;
}
pdata[19] += throughput; pdata[19] += throughput;
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart); } while (!work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce + 1; *hashes_done = pdata[19] - first_nonce;
return 0; return 0;
} }

7
x17/x17.cu

@ -251,11 +251,16 @@ extern "C" int scanhash_x17(int thr_id, struct work* work, uint32_t max_nonce, u
} }
} }
if ((uint64_t)throughput + pdata[19] >= max_nonce) {
pdata[19] = max_nonce;
break;
}
pdata[19] += throughput; pdata[19] += throughput;
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart); } while (pdata[19] < max_nonce && !work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce + 1; *hashes_done = pdata[19] - first_nonce;
return 0; return 0;
} }

Loading…
Cancel
Save