Browse Source

lyra2: small changes and code cleanup

master
Tanguy Pruvot 10 years ago
parent
commit
f9bba0ff1a
  1. 52
      ccminer.cpp
  2. 19
      lyra2/cuda_lyra2.cu
  3. 6
      lyra2/lyra2RE.cu
  4. 4
      miner.h
  5. 2
      util.cpp

52
ccminer.cpp

@ -138,7 +138,7 @@ enum sha_algos { @@ -138,7 +138,7 @@ enum sha_algos {
ALGO_KECCAK,
ALGO_JACKPOT,
ALGO_LUFFA_DOOM,
ALGO_LYRA,
ALGO_LYRA2,
ALGO_MJOLLNIR, /* Hefty hash */
ALGO_MYR_GR,
ALGO_NIST5,
@ -441,25 +441,28 @@ static bool jobj_binary(const json_t *obj, const char *key, @@ -441,25 +441,28 @@ static bool jobj_binary(const json_t *obj, const char *key,
static bool work_decode(const json_t *val, struct work *work)
{
int data_size = sizeof(work->data), target_size = sizeof(work->target);
int adata_sz = ARRAY_SIZE(work->data), atarget_sz = ARRAY_SIZE(work->target);
int i;
if (unlikely(!jobj_binary(val, "data", work->data, sizeof(work->data)))) {
if (unlikely(!jobj_binary(val, "data", work->data, data_size))) {
applog(LOG_ERR, "JSON inval data");
return false;
}
if (unlikely(!jobj_binary(val, "target", work->target, sizeof(work->target)))) {
if (unlikely(!jobj_binary(val, "target", work->target, target_size))) {
applog(LOG_ERR, "JSON inval target");
return false;
}
if (opt_algo == ALGO_HEAVY) {
if (unlikely(!jobj_binary(val, "maxvote", &work->maxvote, sizeof(work->maxvote)))) {
work->maxvote = 2048;
}
} else work->maxvote = 0;
for (i = 0; i < ARRAY_SIZE(work->data); i++)
for (i = 0; i < adata_sz; i++)
work->data[i] = le32dec(work->data + i);
for (i = 0; i < ARRAY_SIZE(work->target); i++)
for (i = 0; i < atarget_sz; i++)
work->target[i] = le32dec(work->target + i);
json_t *jr = json_object_get(val, "noncerange");
@ -468,7 +471,8 @@ static bool work_decode(const json_t *val, struct work *work) @@ -468,7 +471,8 @@ static bool work_decode(const json_t *val, struct work *work)
if (likely(hexstr)) {
// never seen yet...
hex2bin((uchar*)work->noncerange.u64, hexstr, 8);
applog(LOG_DEBUG, "received noncerange: %08x-%08x", work->noncerange.u32[0], work->noncerange.u32[1]);
applog(LOG_DEBUG, "received noncerange: %08x-%08x",
work->noncerange.u32[0], work->noncerange.u32[1]);
}
}
@ -635,7 +639,7 @@ static bool submit_upstream_work(CURL *curl, struct work *work) @@ -635,7 +639,7 @@ static bool submit_upstream_work(CURL *curl, struct work *work)
/* build JSON-RPC request */
sprintf(s,
"{\"method\": \"getwork\", \"params\": [ \"%s\" ], \"id\":1}\r\n",
"{\"method\": \"getwork\", \"params\": [\"%s\"], \"id\":4}\r\n",
str);
/* issue JSON-RPC request */
@ -963,14 +967,23 @@ static void stratum_gen_work(struct stratum_ctx *sctx, struct work *work) @@ -963,14 +967,23 @@ static void stratum_gen_work(struct stratum_ctx *sctx, struct work *work)
free(xnonce2str);
}
if (opt_algo == ALGO_JACKPOT)
diff_to_target(work->target, sctx->job.diff / (65536.0 * opt_difficulty));
else if (opt_algo == ALGO_FUGUE256 || opt_algo == ALGO_GROESTL || opt_algo == ALGO_DMD_GR || opt_algo == ALGO_FRESH)
diff_to_target(work->target, sctx->job.diff / (256.0 * opt_difficulty));
else if (opt_algo == ALGO_KECCAK)
diff_to_target(work->target, sctx->job.diff / (128.0 * opt_difficulty));
else
diff_to_target(work->target, sctx->job.diff / opt_difficulty);
switch (opt_algo) {
case ALGO_JACKPOT:
diff_to_target(work->target, sctx->job.diff / (65536.0 * opt_difficulty));
break;
case ALGO_DMD_GR:
case ALGO_FRESH:
case ALGO_FUGUE256:
case ALGO_GROESTL:
//case ALGO_QUBIT:
diff_to_target(work->target, sctx->job.diff / (256.0 * opt_difficulty));
break;
case ALGO_KECCAK:
diff_to_target(work->target, sctx->job.diff / (128.0 * opt_difficulty));
break;
default:
diff_to_target(work->target, sctx->job.diff / opt_difficulty);
}
}
static void *miner_thread(void *userdata)
@ -1113,6 +1126,9 @@ static void *miner_thread(void *userdata) @@ -1113,6 +1126,9 @@ static void *miner_thread(void *userdata)
case ALGO_X13:
minmax = 0x400000;
break;
case ALGO_LYRA2:
minmax = 0x100000;
break;
}
max64 = max(minmax-1, max64);
}
@ -1263,8 +1279,8 @@ static void *miner_thread(void *userdata) @@ -1263,8 +1279,8 @@ static void *miner_thread(void *userdata)
max_nonce, &hashes_done);
break;
case ALGO_LYRA:
rc = scanhash_lyra(thr_id, work.data, work.target,
case ALGO_LYRA2:
rc = scanhash_lyra2(thr_id, work.data, work.target,
max_nonce, &hashes_done);
break;

19
lyra2/cuda_lyra2.cu

@ -2,6 +2,8 @@ @@ -2,6 +2,8 @@
#include "cuda_helper.h"
#define TPB 256
static __constant__ uint2 blake2b_IV[8] = {
{ 0xf3bcc908, 0x6a09e667 },
{ 0x84caa73b, 0xbb67ae85 },
@ -282,10 +284,9 @@ static __device__ __forceinline__ void round_lyra_v30(uint64_t *s) @@ -282,10 +284,9 @@ static __device__ __forceinline__ void round_lyra_v30(uint64_t *s)
Gfunc_v30(s[3], s[4], s[9], s[14]);
}
__global__ __launch_bounds__(256, 1)
__global__ __launch_bounds__(TPB, 1)
void lyra2_gpu_hash_32_v30(int threads, uint32_t startNounce, uint64_t *outputHash)
{
int thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads)
{
@ -359,10 +360,9 @@ void lyra2_gpu_hash_32_v30(int threads, uint32_t startNounce, uint64_t *outputHa @@ -359,10 +360,9 @@ void lyra2_gpu_hash_32_v30(int threads, uint32_t startNounce, uint64_t *outputHa
} //thread
}
__global__ __launch_bounds__(256, 1)
__global__ __launch_bounds__(TPB, 1)
void lyra2_gpu_hash_32(int threads, uint32_t startNounce, uint64_t *outputHash)
{
int thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads)
{
@ -375,7 +375,7 @@ void lyra2_gpu_hash_32(int threads, uint32_t startNounce, uint64_t *outputHash) @@ -375,7 +375,7 @@ void lyra2_gpu_hash_32(int threads, uint32_t startNounce, uint64_t *outputHash)
for (int i = 0; i<8; i++) { state[i + 8] = blake2b_IV[i]; }
// blake2blyra x2
#pragma unroll 24
//#pragma unroll 24
for (int i = 0; i<24; i++) { round_lyra_v35(state); } //because 12 is not enough
uint2 Matrix[96][8]; // not cool
@ -434,9 +434,9 @@ void lyra2_gpu_hash_32(int threads, uint32_t startNounce, uint64_t *outputHash) @@ -434,9 +434,9 @@ void lyra2_gpu_hash_32(int threads, uint32_t startNounce, uint64_t *outputHash)
} //thread
}
__global__
void __launch_bounds__(256, 1) lyra2_gpu_hash_32_test(int threads, uint32_t startNounce, uint64_t *outputHash)
#if 0
__global__ __launch_bounds__(TPB, 1)
void lyra2_gpu_hash_32_test(int threads, uint32_t startNounce, uint64_t *outputHash)
{
int thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads)
@ -508,6 +508,7 @@ void __launch_bounds__(256, 1) lyra2_gpu_hash_32_test(int threads, uint32_t star @@ -508,6 +508,7 @@ void __launch_bounds__(256, 1) lyra2_gpu_hash_32_test(int threads, uint32_t star
} //thread
}
#endif
__host__
void lyra2_cpu_init(int thr_id, int threads)
@ -518,7 +519,7 @@ void lyra2_cpu_init(int thr_id, int threads) @@ -518,7 +519,7 @@ void lyra2_cpu_init(int thr_id, int threads)
__host__
void lyra2_cpu_hash_32(int thr_id, int threads, uint32_t startNounce, uint64_t *d_outputHash, int order)
{
const int threadsperblock = 256;
const int threadsperblock = TPB;
dim3 grid((threads + threadsperblock - 1) / threadsperblock);
dim3 block(threadsperblock);

6
lyra2/lyra2RE.cu

@ -26,7 +26,7 @@ extern void groestl256_setTarget(const void *ptarget); @@ -26,7 +26,7 @@ extern void groestl256_setTarget(const void *ptarget);
extern uint32_t groestl256_cpu_hash_32(int thr_id, int threads, uint32_t startNounce, uint64_t *d_outputHash, int order);
extern void groestl256_cpu_init(int thr_id, int threads);
extern "C" void lyra_hash(void *state, const void *input)
extern "C" void lyra2_hash(void *state, const void *input)
{
sph_blake256_context ctx_blake;
sph_keccak256_context ctx_keccak;
@ -58,7 +58,7 @@ extern "C" void lyra_hash(void *state, const void *input) @@ -58,7 +58,7 @@ extern "C" void lyra_hash(void *state, const void *input)
static bool init[8] = { 0 };
extern "C" int scanhash_lyra(int thr_id, uint32_t *pdata,
extern "C" int scanhash_lyra2(int thr_id, uint32_t *pdata,
const uint32_t *ptarget, uint32_t max_nonce,
unsigned long *hashes_done)
{
@ -107,7 +107,7 @@ extern "C" int scanhash_lyra(int thr_id, uint32_t *pdata, @@ -107,7 +107,7 @@ extern "C" int scanhash_lyra(int thr_id, uint32_t *pdata,
const uint32_t Htarg = ptarget[7];
uint32_t vhash64[8];
be32enc(&endiandata[19], foundNonce);
lyra_hash(vhash64, endiandata);
lyra2_hash(vhash64, endiandata);
if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) {
*hashes_done = pdata[19] - first_nonce + throughput;

4
miner.h

@ -328,7 +328,7 @@ extern int scanhash_fresh(int thr_id, uint32_t *pdata, @@ -328,7 +328,7 @@ extern int scanhash_fresh(int thr_id, uint32_t *pdata,
const uint32_t *ptarget, uint32_t max_nonce,
unsigned long *hashes_done);
extern int scanhash_lyra(int thr_id, uint32_t *pdata,
extern int scanhash_lyra2(int thr_id, uint32_t *pdata,
const uint32_t *ptarget, uint32_t max_nonce,
unsigned long *hashes_done);
@ -649,7 +649,7 @@ void heavycoin_hash(unsigned char* output, const unsigned char* input, int len); @@ -649,7 +649,7 @@ void heavycoin_hash(unsigned char* output, const unsigned char* input, int len);
void keccak256_hash(void *state, const void *input);
unsigned int jackpothash(void *state, const void *input);
void groestlhash(void *state, const void *input);
void lyra_hash(void *state, const void *input);
void lyra2_hash(void *state, const void *input);
void myriadhash(void *state, const void *input);
void nist5hash(void *state, const void *input);
void pentablakehash(void *output, const void *input);

2
util.cpp

@ -1646,7 +1646,7 @@ void print_hash_tests(void) @@ -1646,7 +1646,7 @@ void print_hash_tests(void)
printpfx("luffa", hash);
memset(hash, 0, sizeof hash);
lyra_hash(&hash[0], &buf[0]);
lyra2_hash(&hash[0], &buf[0]);
printpfx("lyra2", hash);
memset(hash, 0, sizeof hash);

Loading…
Cancel
Save