GOSTSec
/
ccminer-gostd-lite
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity
Browse Source

lyra2v2: increase default intensity 

to be able to say, like sp, that its faster :p
master
Tanguy Pruvot 9 years ago
parent
b3adebdf2a
commit
87edf84bf3
2 changed files with 16 additions and 16 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 4
      lyra2/cuda_lyra2v2.cu
  2. 28
      lyra2/lyra2REv2.cu

4
lyra2/cuda_lyra2v2.cu
Unescape View File

@ -350,10 +350,10 @@ __global__ void lyra2v2_gpu_hash_32(uint32_t threads, uint32_t startNounce, uint
#endif #endif
__host__ __host__
void lyra2v2_cpu_init(int thr_id, uint32_t threads, uint64_t *d_hash2) void lyra2v2_cpu_init(int thr_id, uint32_t threads, uint64_t *d_matrix)
{ {
// just assign the device pointer allocated in main loop // just assign the device pointer allocated in main loop
cudaMemcpyToSymbol(DMatrix, &d_hash2, sizeof(uint64_t*), 0, cudaMemcpyHostToDevice); cudaMemcpyToSymbol(DMatrix, &d_matrix, sizeof(uint64_t*), 0, cudaMemcpyHostToDevice);
} }
__host__ __host__

28
lyra2/lyra2REv2.cu
Unescape View File

@ -12,7 +12,7 @@ extern "C" {
static _ALIGN(64) uint64_t *d_hash[MAX_GPUS]; static _ALIGN(64) uint64_t *d_hash[MAX_GPUS];
static uint64_t *d_hash2[MAX_GPUS]; static uint64_t* d_matrix[MAX_GPUS];
extern void blake256_cpu_init(int thr_id, uint32_t threads); extern void blake256_cpu_init(int thr_id, uint32_t threads);
extern void blake256_cpu_hash_80(const int thr_id, const uint32_t threads, const uint32_t startNonce, uint64_t *Hash, int order); extern void blake256_cpu_hash_80(const int thr_id, const uint32_t threads, const uint32_t startNonce, uint64_t *Hash, int order);
@ -24,7 +24,7 @@ extern void skein256_cpu_init(int thr_id, uint32_t threads);
extern void cubehash256_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startNounce, uint64_t *d_hash, int order); extern void cubehash256_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startNounce, uint64_t *d_hash, int order);
extern void lyra2v2_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startNonce, uint64_t *d_outputHash, int order); extern void lyra2v2_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startNonce, uint64_t *d_outputHash, int order);
extern void lyra2v2_cpu_init(int thr_id, uint32_t threads, uint64_t* matrix); extern void lyra2v2_cpu_init(int thr_id, uint32_t threads, uint64_t* d_matrix);
extern void bmw256_setTarget(const void *ptarget); extern void bmw256_setTarget(const void *ptarget);
extern void bmw256_cpu_init(int thr_id, uint32_t threads); extern void bmw256_cpu_init(int thr_id, uint32_t threads);
@ -77,9 +77,10 @@ extern "C" int scanhash_lyra2v2(int thr_id, struct work* work, uint32_t max_nonc
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];
int intensity = (device_sm[device_map[thr_id]] > 500 && !is_windows()) ? 18 : 17; int dev_id = device_map[thr_id];
int intensity = (device_sm[dev_id] > 500 && !is_windows()) ? 20 : 18;
unsigned int defthr = 1U << intensity; unsigned int defthr = 1U << intensity;
uint32_t throughput = device_intensity(device_map[thr_id], __func__, defthr); uint32_t throughput = device_intensity(dev_id, __func__, defthr);
if (opt_benchmark) if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x00ff; ((uint32_t*)ptarget)[7] = 0x00ff;
@ -102,10 +103,10 @@ extern "C" int scanhash_lyra2v2(int thr_id, struct work* work, uint32_t max_nonc
} }
// DMatrix // DMatrix
CUDA_SAFE_CALL(cudaMalloc(&d_hash2[thr_id], 16 * 4 * 4 * sizeof(uint64_t) * throughput)); CUDA_SAFE_CALL(cudaMalloc(&d_matrix[thr_id], (size_t)16 * sizeof(uint64_t) * 4 * 3 * throughput));
lyra2v2_cpu_init(thr_id, throughput, d_hash2[thr_id]); lyra2v2_cpu_init(thr_id, throughput, d_matrix[thr_id]);
CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], (size_t)throughput * 32)); CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], (size_t)32 * throughput));
init[thr_id] = true; init[thr_id] = true;
} }
@ -130,18 +131,18 @@ extern "C" int scanhash_lyra2v2(int thr_id, struct work* work, uint32_t max_nonc
bmw256_cpu_hash_32(thr_id, throughput, pdata[19], d_hash[thr_id], foundNonces); bmw256_cpu_hash_32(thr_id, throughput, pdata[19], d_hash[thr_id], foundNonces);
*hashes_done = pdata[19] - first_nonce + throughput;
if (foundNonces[0] != 0) if (foundNonces[0] != 0)
{ {
const uint32_t Htarg = ptarget[7];
uint32_t vhash64[8]; uint32_t vhash64[8];
be32enc(&endiandata[19], foundNonces[0]); be32enc(&endiandata[19], foundNonces[0]);
lyra2v2_hash(vhash64, endiandata); lyra2v2_hash(vhash64, endiandata);
if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) if (vhash64[7] <= ptarget[7] && fulltest(vhash64, ptarget))
{ {
int res = 1; int res = 1;
bn_store_hash_target_ratio(vhash64, ptarget, work); bn_store_hash_target_ratio(vhash64, ptarget, work);
// check if there was another one... // check if there was another one...
*hashes_done = pdata[19] - first_nonce + throughput;
if (foundNonces[1] != 0) if (foundNonces[1] != 0)
{ {
be32enc(&endiandata[19], foundNonces[1]); be32enc(&endiandata[19], foundNonces[1]);
@ -157,8 +158,7 @@ extern "C" int scanhash_lyra2v2(int thr_id, struct work* work, uint32_t max_nonc
} }
else else
{ {
if (vhash64[7] > Htarg) // don't show message if it is equal but fails fulltest applog(LOG_WARNING, "GPU #%d: result does not validate on CPU!", dev_id);
applog(LOG_WARNING, "GPU #%d: result does not validate on CPU!", thr_id);
} }
} }
@ -180,10 +180,10 @@ extern "C" void free_lyra2v2(int thr_id)
cudaSetDevice(device_map[thr_id]); cudaSetDevice(device_map[thr_id]);
cudaFree(d_hash[thr_id]); cudaFree(d_hash[thr_id]);
cudaFree(d_hash2[thr_id]); cudaFree(d_matrix[thr_id]);
bmw256_cpu_free(thr_id); bmw256_cpu_free(thr_id);
init[thr_id] = false; init[thr_id] = false;
cudaDeviceSynchronize(); cudaDeviceSynchronize();
} }

Write Preview
Loading…
Cancel
Save