Browse Source

make throughput unsigned

Note: The next thing to do is to change all "int threads" in cuda
KlausT 10 years ago committed by Tanguy Pruvot
parent
commit
c6baddc85a
  1. 4
      JHA/jackpotcoin.cu
  2. 2
      cuda_nist5.cu
  3. 6
      heavy/heavy.cu
  4. 4
      lyra2/lyra2RE.cu
  5. 5
      pentablake.cu
  6. 5
      quark/quarkcoin.cu
  7. 2
      qubit/deep.cu
  8. 5
      qubit/qubit.cu
  9. 11
      x11/s3.cu
  10. 5
      x11/x11.cu
  11. 4
      x15/whirlpool.cu
  12. 4
      x15/x14.cu
  13. 4
      x15/x15.cu
  14. 4
      x17/x17.cu

4
JHA/jackpotcoin.cu

@ -96,8 +96,8 @@ extern "C" int scanhash_jackpot(int thr_id, uint32_t *pdata,
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x000f; ((uint32_t*)ptarget)[7] = 0x000f;
int throughput = opt_work_size ? opt_work_size : (1 << 20); // 256*4096 uint32_t throughput = opt_work_size ? opt_work_size : (1 << 20); // 256*4096
throughput = min(throughput, (int)(max_nonce - first_nonce)); throughput = min(throughput, max_nonce - first_nonce);
if (!init[thr_id]) if (!init[thr_id])
{ {

2
cuda_nist5.cu

@ -75,7 +75,7 @@ extern "C" int scanhash_nist5(int thr_id, uint32_t *pdata,
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x00FF; ((uint32_t*)ptarget)[7] = 0x00FF;
int throughput = opt_work_size ? opt_work_size : (1 << 20); // 256*4096 uint32_t throughput = opt_work_size ? opt_work_size : (1 << 20); // 256*4096
throughput = min(throughput, (int) (max_nonce - first_nonce)); throughput = min(throughput, (int) (max_nonce - first_nonce));
if (!init[thr_id]) if (!init[thr_id])

6
heavy/heavy.cu

@ -127,7 +127,7 @@ struct check_nonce_for_remove
} }
}; };
static bool init[8] = {0,0,0,0,0,0,0,0}; static bool init[8] = { 0 };
__host__ __host__
int scanhash_heavy(int thr_id, uint32_t *pdata, int scanhash_heavy(int thr_id, uint32_t *pdata,
@ -135,8 +135,8 @@ int scanhash_heavy(int thr_id, uint32_t *pdata,
unsigned long *hashes_done, uint32_t maxvote, int blocklen) unsigned long *hashes_done, uint32_t maxvote, int blocklen)
{ {
const uint32_t first_nonce = pdata[19]; const uint32_t first_nonce = pdata[19];
// CUDA will process thousands of threads. // Remove 256 to allow -i 20
int throughput = opt_work_size ? opt_work_size : (1 << 19) - 256; // 256*2048 uint32_t throughput = opt_work_size ? opt_work_size : (1 << 19) - 256; // 256*2048
throughput = min(throughput, (int)(max_nonce - first_nonce)); throughput = min(throughput, (int)(max_nonce - first_nonce));
int rc = 0; int rc = 0;

4
lyra2/lyra2RE.cu

@ -64,8 +64,8 @@ extern "C" int scanhash_lyra2(int thr_id, uint32_t *pdata,
{ {
const uint32_t first_nonce = pdata[19]; const uint32_t first_nonce = pdata[19];
int intensity = (device_sm[device_map[thr_id]] >= 500 && !is_windows()) ? 18 : 17; int intensity = (device_sm[device_map[thr_id]] >= 500 && !is_windows()) ? 18 : 17;
int throughput = opt_work_size ? opt_work_size : (1 << intensity); // 18=256*256*4; uint32_t throughput = opt_work_size ? opt_work_size : (1 << intensity); // 18=256*256*4;
throughput = min(throughput, (int)(max_nonce - first_nonce)); throughput = min(throughput, max_nonce - first_nonce);
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x0000ff; ((uint32_t*)ptarget)[7] = 0x0000ff;

5
pentablake.cu

@ -492,8 +492,9 @@ extern "C" int scanhash_pentablake(int thr_id, uint32_t *pdata, const uint32_t *
const uint32_t first_nonce = pdata[19]; const uint32_t first_nonce = pdata[19];
uint32_t endiandata[20]; uint32_t endiandata[20];
int rc = 0; int rc = 0;
int throughput = opt_work_size ? opt_work_size : (128 * 2560); // 18.5 uint32_t throughput = opt_work_size ? opt_work_size : (128 * 2560); // 18.5
throughput = min(throughput, (int)(max_nonce - first_nonce));
throughput = min(throughput, max_nonce - first_nonce);
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x000F; ((uint32_t*)ptarget)[7] = 0x000F;

5
quark/quarkcoin.cu

@ -137,9 +137,8 @@ extern "C" int scanhash_quark(int thr_id, uint32_t *pdata,
unsigned long *hashes_done) unsigned long *hashes_done)
{ {
const uint32_t first_nonce = pdata[19]; const uint32_t first_nonce = pdata[19];
uint32_t throughput = opt_work_size ? opt_work_size : (1 << 20); // 256*4096
int throughput = opt_work_size ? opt_work_size : (1 << 20); // 256*4096 throughput = min(throughput, max_nonce - first_nonce);
throughput = min(throughput, (int)(max_nonce - first_nonce));
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x00F; ((uint32_t*)ptarget)[7] = 0x00F;

2
qubit/deep.cu

@ -69,7 +69,7 @@ extern "C" int scanhash_deep(int thr_id, uint32_t *pdata,
if (!init[thr_id]) if (!init[thr_id])
{ {
cudaSetDevice(device_map[thr_id]); cudaSetDevice(device_map[thr_id]);
cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput); CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput));
qubit_luffa512_cpu_init(thr_id, throughput); qubit_luffa512_cpu_init(thr_id, throughput);
x11_cubehash512_cpu_init(thr_id, throughput); x11_cubehash512_cpu_init(thr_id, throughput);

5
qubit/qubit.cu

@ -80,8 +80,9 @@ extern "C" int scanhash_qubit(int thr_id, uint32_t *pdata,
{ {
uint32_t endiandata[20]; uint32_t endiandata[20];
const uint32_t first_nonce = pdata[19]; const uint32_t first_nonce = pdata[19];
int throughput = opt_work_size ? opt_work_size : (1 << 19); // 256*256*8 uint32_t throughput = opt_work_size ? opt_work_size : (1 << 19); // 256*256*8
throughput = min(throughput, (int)(max_nonce - first_nonce));
throughput = min(throughput, max_nonce - first_nonce);
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x0000ff; ((uint32_t*)ptarget)[7] = 0x0000ff;

11
x11/s3.cu

@ -57,13 +57,10 @@ extern "C" int scanhash_s3(int thr_id, uint32_t *pdata,
unsigned long *hashes_done) unsigned long *hashes_done)
{ {
const uint32_t first_nonce = pdata[19]; const uint32_t first_nonce = pdata[19];
int intensity = 20; // 256*256*8*2; int intensity = (device_sm[device_map[thr_id]] >= 500 && !is_windows()) ? 20 : 19;
#ifdef WIN32 uint32_t throughput = opt_work_size ? opt_work_size : (1 << intensity);
// reduce by one the intensity on windows
intensity--; throughput = min(throughput, max_nonce - first_nonce);
#endif
int throughput = opt_work_size ? opt_work_size : (1 << intensity);
throughput = min(throughput, (int)(max_nonce - first_nonce));
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0xF; ((uint32_t*)ptarget)[7] = 0xF;

5
x11/x11.cu

@ -133,8 +133,9 @@ extern "C" int scanhash_x11(int thr_id, uint32_t *pdata,
{ {
const uint32_t first_nonce = pdata[19]; const uint32_t first_nonce = pdata[19];
int intensity = (device_sm[device_map[thr_id]] >= 500 && !is_windows()) ? 20 : 19; int intensity = (device_sm[device_map[thr_id]] >= 500 && !is_windows()) ? 20 : 19;
int throughput = opt_work_size ? opt_work_size : (1 << intensity); // 20=256*256*16; uint32_t throughput = opt_work_size ? opt_work_size : intensity; // 20=256*256*16;
throughput = min(throughput, (int)(max_nonce - first_nonce));
throughput = min(throughput, max_nonce - first_nonce);
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x5; ((uint32_t*)ptarget)[7] = 0x5;

4
x15/whirlpool.cu

@ -57,8 +57,8 @@ extern "C" int scanhash_whc(int thr_id, uint32_t *pdata,
{ {
const uint32_t first_nonce = pdata[19]; const uint32_t first_nonce = pdata[19];
uint32_t endiandata[20]; uint32_t endiandata[20];
int throughput = opt_work_size ? opt_work_size : (1 << 19); // 256*256*8; uint32_t throughput = opt_work_size ? opt_work_size : (1 << 19); // 256*256*8;
throughput = min(throughput, (int)(max_nonce - first_nonce)); throughput = min(throughput, max_nonce - first_nonce);
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x0000ff; ((uint32_t*)ptarget)[7] = 0x0000ff;

4
x15/x14.cu

@ -163,8 +163,8 @@ extern "C" int scanhash_x14(int thr_id, uint32_t *pdata,
{ {
const uint32_t first_nonce = pdata[19]; const uint32_t first_nonce = pdata[19];
uint32_t endiandata[20]; uint32_t endiandata[20];
int throughput = opt_work_size ? opt_work_size : (1 << 19); // 256*256*8; uint32_t throughput = opt_work_size ? opt_work_size : (1 << 19); // 256*256*8;
throughput = min(throughput, (int)(max_nonce - first_nonce)); throughput = min(throughput, max_nonce - first_nonce);
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x000f; ((uint32_t*)ptarget)[7] = 0x000f;

4
x15/x15.cu

@ -174,8 +174,8 @@ extern "C" int scanhash_x15(int thr_id, uint32_t *pdata,
const uint32_t first_nonce = pdata[19]; const uint32_t first_nonce = pdata[19];
uint32_t endiandata[20]; uint32_t endiandata[20];
int throughput = opt_work_size ? opt_work_size : (1 << 19); // 256*256*8; uint32_t throughput = opt_work_size ? opt_work_size : (1 << 19); // 256*256*8;
throughput = min(throughput, (int)(max_nonce - first_nonce)); throughput = min(throughput, max_nonce - first_nonce);
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x00FF; ((uint32_t*)ptarget)[7] = 0x00FF;

4
x17/x17.cu

@ -191,8 +191,8 @@ extern "C" int scanhash_x17(int thr_id, uint32_t *pdata,
unsigned long *hashes_done) unsigned long *hashes_done)
{ {
const uint32_t first_nonce = pdata[19]; const uint32_t first_nonce = pdata[19];
int throughput = opt_work_size ? opt_work_size : (1 << 19); // 256*256*8; uint32_t throughput = opt_work_size ? opt_work_size : (1 << 19); // 256*256*8;
throughput = min(throughput, (int)(max_nonce - first_nonce)); throughput = min(throughput, max_nonce - first_nonce);
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x00ff; ((uint32_t*)ptarget)[7] = 0x00ff;

Loading…
Cancel
Save