diff --git a/ccminer.cpp b/ccminer.cpp index 4cc0d46..7392886 100644 --- a/ccminer.cpp +++ b/ccminer.cpp @@ -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, 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) 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) /* 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) 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) 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) 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; diff --git a/lyra2/cuda_lyra2.cu b/lyra2/cuda_lyra2.cu index 3508d34..95d61f5 100644 --- a/lyra2/cuda_lyra2.cu +++ b/lyra2/cuda_lyra2.cu @@ -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) 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 } //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) 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) } //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 } //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) __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); diff --git a/lyra2/lyra2RE.cu b/lyra2/lyra2RE.cu index 60ed22a..d1b047f 100644 --- a/lyra2/lyra2RE.cu +++ b/lyra2/lyra2RE.cu @@ -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) 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, 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; diff --git a/miner.h b/miner.h index 59e3423..0a4219e 100644 --- a/miner.h +++ b/miner.h @@ -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); 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); diff --git a/util.cpp b/util.cpp index 07564f4..e88f9df 100644 --- a/util.cpp +++ b/util.cpp @@ -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);