From a65ab2ae18b769fd7d91fb13f770c2fa1d87fe39 Mon Sep 17 00:00:00 2001 From: iwasaki Date: Sat, 5 Jan 2019 17:52:25 +0900 Subject: [PATCH] sapling --- ccminer.cpp | 104 +++++++++++++++++++++++++++----- miner.h | 4 +- util.cpp | 19 +++++- yescrypt/cuda_yescrypt.cu | 123 +++++++++++++++++++++++++++++++++++--- yescrypt/yescrypt.cu | 55 +++++++++-------- 5 files changed, 251 insertions(+), 54 deletions(-) diff --git a/ccminer.cpp b/ccminer.cpp index 6d8ad73..8e2e981 100644 --- a/ccminer.cpp +++ b/ccminer.cpp @@ -831,24 +831,25 @@ static bool submit_upstream_work(CURL *curl, struct work *work) char data_str[2 * sizeof(work->data) + 1]; char *req; + int datasize = work->sapling ? 112 : 80; for (int i = 0; i < ARRAY_SIZE(work->data); i++) be32enc(work->data + i, work->data[i]); - cbin2hex(data_str, (char *)work->data, 80); + cbin2hex(data_str, (char *)work->data, datasize); if (work->workid) { char *params; val = json_object(); json_object_set_new(val, "workid", json_string(work->workid)); params = json_dumps(val, 0); json_decref(val); - req = (char*)malloc(128 + 2 * 80 + strlen(work->txs2) + strlen(params)); + req = (char*)malloc(128 + 2 * datasize + strlen(work->txs2) + strlen(params)); sprintf(req, "{\"method\": \"submitblock\", \"params\": [\"%s%s\", %s], \"id\":4}\r\n", data_str, work->txs2, params); free(params); } else { - req = (char*)malloc(128 + 2 * 80 + strlen(work->txs2)); + req = (char*)malloc(128 + 2 * 2 * datasize + strlen(work->txs2)); sprintf(req, "{\"method\": \"submitblock\", \"params\": [\"%s%s\"], \"id\":4}\r\n", data_str, work->txs2); @@ -862,8 +863,27 @@ static bool submit_upstream_work(CURL *curl, struct work *work) } res = json_object_get(val, "result"); - reason = json_object_get(val, "reject-reason"); - if (!share_result(json_is_true(res), reason ? json_string_value(reason) : NULL)) + int ret; + if (json_is_object(res)) + { + char *res_str; + bool sumres = false; + void *iter = json_object_iter(res); + while (iter) { + if (json_is_null(json_object_iter_value(iter)))\ + { + sumres = true; + break; + } + iter = json_object_iter_next(res, iter); + } + res_str = json_dumps(res, 0); + ret = share_result(sumres, res_str); + free(res_str); + } else { + ret = share_result(json_is_null(res), json_string_value(res)); + } + if (!ret) { if (check_dups) hashlog_purge_job(work->job_id); @@ -875,7 +895,6 @@ static bool submit_upstream_work(CURL *curl, struct work *work) #endif else { - /* build hex string */ char *str = NULL; for(int i = 0; i < (work->datasize >> 2); i++) @@ -901,8 +920,27 @@ static bool submit_upstream_work(CURL *curl, struct work *work) } res = json_object_get(val, "result"); - reason = json_object_get(val, "reject-reason"); - if(!share_result(json_is_true(res), reason ? json_string_value(reason) : NULL)) + int ret; + if (json_is_object(res)) + { + char *res_str; + bool sumres = false; + void *iter = json_object_iter(res); + while (iter) { + if (json_is_null(json_object_iter_value(iter)))\ + { + sumres = true; + break; + } + iter = json_object_iter_next(res, iter); + } + res_str = json_dumps(res, 0); + ret = share_result(sumres, res_str); + free(res_str); + } else { + ret = share_result(json_is_null(res), json_string_value(res)); + } + if (!ret) { if(check_dups) hashlog_purge_job(work->job_id); @@ -1015,12 +1053,14 @@ static bool gbt_work_decode_full(const json_t *val, struct work *work) int tx_count, tx_size; uchar txc_vi[9]; uchar(*merkle_tree)[32] = NULL; + uint32_t final_sapling_hash[8]; bool coinbase_append = false; bool submit_coinbase = false; bool version_force = false; bool version_reduce = false; json_t *tmp, *txa; bool rc = false; + bool sapling = false; tmp = json_object_get(val, "mutable"); if (tmp && json_is_array(tmp)) { @@ -1054,7 +1094,9 @@ static bool gbt_work_decode_full(const json_t *val, struct work *work) goto out; } version = (uint32_t)json_integer_value(tmp); - if ((version & 0xffU) > BLOCK_VERSION_CURRENT) { + if ((version & 0xffU) == 5) { + sapling = true; + } else if ((version & 0xffU) > BLOCK_VERSION_CURRENT) { if (version_reduce) { version = (version & ~0xffU) | BLOCK_VERSION_CURRENT; } @@ -1086,6 +1128,13 @@ static bool gbt_work_decode_full(const json_t *val, struct work *work) goto out; } + if (sapling) { + if (unlikely(!jobj_binary(val, "finalsaplingroothash", final_sapling_hash, sizeof(final_sapling_hash)))) { + applog(LOG_ERR, "JSON invalid finalsaplingroothash"); + goto out; + } + } + /* find count and size of transactions */ txa = json_object_get(val, "transactions"); if (!txa || !json_is_array(txa)) { @@ -1236,9 +1285,21 @@ static bool gbt_work_decode_full(const json_t *val, struct work *work) work->data[9 + i] = be32dec((uint32_t *)merkle_tree[0] + i); work->data[17] = swab32(curtime); work->data[18] = le32dec(&bits); - memset(work->data + 19, 0x00, 52); - work->data[20] = 0x80000000; - work->data[31] = 0x00000280; + if (sapling) { + work->sapling = true; + work->data[19] = 0x00000000; + for (i = 0; i < 8; i++) + work->data[27 - i] = le32dec(final_sapling_hash + i); + work->data[28] = 0x80000000; + work->data[29] = 0x00000000; + work->data[30] = 0x00000000; + work->data[31] = 0x00000380; + } else { + work->sapling = false; + memset(work->data + 19, 0x00, 52); + work->data[20] = 0x80000000; + work->data[31] = 0x00000280; + } if (unlikely(!jobj_binary(val, "target", target, sizeof(target)))) { applog(LOG_ERR, "JSON invalid target"); @@ -1790,8 +1851,19 @@ static bool stratum_gen_work(struct stratum_ctx *sctx, struct work *work) work->data[9 + i] = be32dec((uint32_t *)merkle_root + i); work->data[17] = le32dec(sctx->job.ntime); work->data[18] = le32dec(sctx->job.nbits); - work->data[20] = 0x80000000; - work->data[31] = 0x00000280; + if (be32dec(sctx->job.version) >= 5) { + work->sapling = true; + for (i = 0; i < 8; i++) + work->data[20 + i] = le32dec((uint32_t *)sctx->job.finalsaplinghash + i); + work->data[28] = 0x80000000; + work->data[29] = 0x00000000; + work->data[30] = 0x00000000; + work->data[31] = 0x00000380; + } else { + work->sapling = false; + work->data[20] = 0x80000000; + work->data[31] = 0x00000280; + } } else { @@ -1871,6 +1943,7 @@ static void *miner_thread(void *userdata) bool extrajob = false; char s[16]; int rc = 0; + int perslen; memset(&work, 0, sizeof(work)); // prevent work from being used uninitialized @@ -2302,8 +2375,9 @@ static void *miner_thread(void *userdata) #ifndef ORG case ALGO_YESCRYPT: + perslen = work.sapling ? 112 : 80; rc = scanhash_yescrypt(thr_id, work.data, work.target, - max_nonce, &hashes_done); + max_nonce, &hashes_done, perslen); break; case ALGO_YESCRYPTR8: diff --git a/miner.h b/miner.h index 12de82a..b57450a 100644 --- a/miner.h +++ b/miner.h @@ -379,7 +379,7 @@ extern int scanhash_x17(int thr_id, uint32_t *pdata, #ifndef ORG extern int scanhash_yescrypt(int thr_id, uint32_t *pdata, uint32_t *ptarget, uint32_t max_nonce, - uint32_t *hashes_done); + uint32_t *hashes_done, int perslen); extern int scanhash_yescryptr8(int thr_id, uint32_t *pdata, uint32_t *ptarget, uint32_t max_nonce, @@ -612,6 +612,7 @@ uint32_t device_intensity(int thr_id, const char *func, uint32_t defcount); struct stratum_job { char *job_id; unsigned char prevhash[32]; + unsigned char finalsaplinghash[32]; size_t coinbase_size; unsigned char *coinbase; unsigned char *xnonce2; @@ -677,6 +678,7 @@ struct work { char *txs2; char *workid; #endif + bool sapling; }; enum sha_algos diff --git a/util.cpp b/util.cpp index 1cab00f..ebcc49d 100644 --- a/util.cpp +++ b/util.cpp @@ -146,7 +146,7 @@ void applog(int prio, const char *fmt, ...) { case LOG_ERR: color = CL_RED; break; case LOG_WARNING: color = CL_YLW; break; - case LOG_NOTICE: color = CL_WHT; break; + case LOG_NOTICE: color = CL_LGR; break; case LOG_INFO: color = ""; break; case LOG_DEBUG: color = CL_GRY; break; @@ -1526,7 +1526,7 @@ static uint32_t getblocheight(struct stratum_ctx *sctx) static bool stratum_notify(struct stratum_ctx *sctx, json_t *params) { - const char *job_id, *prevhash, *coinb1, *coinb2, *version, *nbits, *nreward; + const char *job_id, *prevhash, *coinb1, *coinb2, *version, *nbits, *nreward, *finalsaplinghash; char *stime; size_t coinb1_size, coinb2_size; bool clean, ret = false; @@ -1534,6 +1534,7 @@ static bool stratum_notify(struct stratum_ctx *sctx, json_t *params) json_t *merkle_arr; uchar **merkle = NULL; int32_t ntime; + int ver; job_id = json_string_value(json_array_get(params, 0)); prevhash = json_string_value(json_array_get(params, 1)); @@ -1575,6 +1576,16 @@ static bool stratum_notify(struct stratum_ctx *sctx, json_t *params) } } + hex2bin(sctx->job.version, version, 4); + ver = be32dec(sctx->job.version); + if (ver == 5) { + finalsaplinghash = json_string_value(json_array_get(params, 9)); + if (!finalsaplinghash || strlen(finalsaplinghash) != 64) { + applog(LOG_ERR, "Stratum notify: invalid parameters"); + goto out; + } + } + /* store stratum server time diff */ hex2bin((uchar *)&ntime, stime, 4); if(opt_algo!=ALGO_SIA) @@ -1643,6 +1654,9 @@ static bool stratum_notify(struct stratum_ctx *sctx, json_t *params) free(sctx->job.job_id); sctx->job.job_id = strdup(job_id); hex2bin(sctx->job.prevhash, prevhash, 32); + if (ver == 5) { + hex2bin(sctx->job.finalsaplinghash, finalsaplinghash, 32); + } if(opt_algo != ALGO_SIA) sctx->job.height = getblocheight(sctx); @@ -1659,7 +1673,6 @@ static bool stratum_notify(struct stratum_ctx *sctx, json_t *params) sctx->job.merkle = merkle; sctx->job.merkle_count = merkle_count; - hex2bin(sctx->job.version, version, 4); hex2bin(sctx->job.nbits, nbits, 4); hex2bin(sctx->job.ntime, stime, 4); if(nreward != NULL) diff --git a/yescrypt/cuda_yescrypt.cu b/yescrypt/cuda_yescrypt.cu index 6a17841..5a0c4fd 100644 --- a/yescrypt/cuda_yescrypt.cu +++ b/yescrypt/cuda_yescrypt.cu @@ -152,7 +152,7 @@ static uint32_t *d_GNonce[MAX_GPUS]; /////////////////////////////////////////////////////////////////////////////////// /////////////////////////////// sha256 function /////////////////////////////////// -__constant__ static uint32_t c_data[20]; +__constant__ static uint32_t c_data[28]; __constant__ static uint32_t cpu_h[8]; __constant__ static uint32_t c_K[64]; __constant__ static uint32_t client_key[32]; @@ -502,6 +502,101 @@ __global__ void yescrypt_gpu_hash_k0(int threads, uint32_t startNonce, const uin } } +__launch_bounds__(32, 1) +__global__ void yescrypt_gpu_hash_k0_112bytes(int threads, uint32_t startNonce, const uint32_t r, const uint32_t p) +{ + int thread = (blockDim.x * blockIdx.x + threadIdx.x); + + //if (thread < threads) + { + uint32_t nonce = startNonce + thread; + uint32_t in[16]; + uint32_t result[16]; + uint32_t state1[8], state2[8]; + uint32_t passwd[8]; + + in[0] = c_data[16]; in[1] = c_data[17]; in[2] = c_data[18]; in[3] = nonce; + in[4] = c_data[20]; in[5] = c_data[21]; in[6] = c_data[22]; in[7] = c_data[23]; + in[8] = c_data[24]; in[9] = c_data[25]; in[10] = c_data[26]; in[11] = c_data[27]; + in[12] = 0x80000000; in[13] = in[14] = 0x00000000; in[15] = 0x00000380; + passwd[0] = cpu_h[0]; passwd[1] = cpu_h[1]; passwd[2] = cpu_h[2]; passwd[3] = cpu_h[3]; + passwd[4] = cpu_h[4]; passwd[5] = cpu_h[5]; passwd[6] = cpu_h[6]; passwd[7] = cpu_h[7]; + sha256_round_body(in, passwd); // length = 112 * 8 = 896 = 0x380 + + in[0] = passwd[0] ^ 0x36363636; in[1] = passwd[1] ^ 0x36363636; in[2] = passwd[2] ^ 0x36363636; in[3] = passwd[3] ^ 0x36363636; + in[4] = passwd[4] ^ 0x36363636; in[5] = passwd[5] ^ 0x36363636; in[6] = passwd[6] ^ 0x36363636; in[7] = passwd[7] ^ 0x36363636; + in[8] = in[9] = in[10] = in[11] = in[12] = in[13] = in[14] = in[15] = 0x36363636; + state1[0] = 0x6A09E667; state1[1] = 0xBB67AE85; state1[2] = 0x3C6EF372; state1[3] = 0xA54FF53A; + state1[4] = 0x510E527F; state1[5] = 0x9B05688C; state1[6] = 0x1F83D9AB; state1[7] = 0x5BE0CD19; + sha256_round_body(in, state1); // inner 64byte + + in[0] = passwd[0] ^ 0x5c5c5c5c; in[1] = passwd[1] ^ 0x5c5c5c5c; in[2] = passwd[2] ^ 0x5c5c5c5c; in[3] = passwd[3] ^ 0x5c5c5c5c; + in[4] = passwd[4] ^ 0x5c5c5c5c; in[5] = passwd[5] ^ 0x5c5c5c5c; in[6] = passwd[6] ^ 0x5c5c5c5c; in[7] = passwd[7] ^ 0x5c5c5c5c; + in[8] = in[9] = in[10] = in[11] = in[12] = in[13] = in[14] = in[15] = 0x5c5c5c5c; + state2[0] = 0x6A09E667; state2[1] = 0xBB67AE85; state2[2] = 0x3C6EF372; state2[3] = 0xA54FF53A; + state2[4] = 0x510E527F; state2[5] = 0x9B05688C; state2[6] = 0x1F83D9AB; state2[7] = 0x5BE0CD19; + sha256_round_body(in, state2); // outer 64byte + + in[0] = c_data[0]; in[1] = c_data[1]; in[2] = c_data[2]; in[3] = c_data[3]; + in[4] = c_data[4]; in[5] = c_data[5]; in[6] = c_data[6]; in[7] = c_data[7]; + in[8] = c_data[8]; in[9] = c_data[9]; in[10] = c_data[10]; in[11] = c_data[11]; + in[12] = c_data[12]; in[13] = c_data[13]; in[14] = c_data[14]; in[15] = c_data[15]; + sha256_round_body(in, state1); // inner 128byte + +#pragma unroll + for (uint32_t i = 0; i < 2 * r*p; i++) + { + in[0] = c_data[16]; in[1] = c_data[17]; in[2] = c_data[18]; in[3] = nonce; + in[4] = c_data[20]; in[5] = c_data[21]; in[6] = c_data[22]; in[7] = c_data[23]; + in[8] = c_data[24]; in[9] = c_data[25]; in[10] = c_data[26]; in[11] = c_data[27]; + in[12] = i * 2 + 1; in[13] = 0x80000000; in[14] = 0x00000000; in[15] = 0x000005A0; + result[0] = state1[0]; result[1] = state1[1]; result[2] = state1[2]; result[3] = state1[3]; + result[4] = state1[4]; result[5] = state1[5]; result[6] = state1[6]; result[7] = state1[7]; + sha256_round_body(in, result + 0); // inner length = 180 * 8 = 1184 = 0x5A0 + + in[0] = result[0]; in[1] = result[1]; in[2] = result[2]; in[3] = result[3]; + in[4] = result[4]; in[5] = result[5]; in[6] = result[6]; in[7] = result[7]; + in[8] = 0x80000000; in[15] = 0x00000300; + in[9] = in[10] = in[11] = in[12] = in[13] = in[14] = 0x00000000; + result[0] = state2[0]; result[1] = state2[1]; result[2] = state2[2]; result[3] = state2[3]; + result[4] = state2[4]; result[5] = state2[5]; result[6] = state2[6]; result[7] = state2[7]; + sha256_round_body(in, result + 0); // outer length = 96 * 8 = 768 = 0x300 + + in[0] = c_data[16]; in[1] = c_data[17]; in[2] = c_data[18]; in[3] = nonce; + in[4] = c_data[20]; in[5] = c_data[21]; in[6] = c_data[22]; in[7] = c_data[23]; + in[8] = c_data[24]; in[9] = c_data[25]; in[10] = c_data[26]; in[11] = c_data[27]; + in[12] = i * 2 + 2; in[13] = 0x80000000; in[14] = 0x00000000; in[15] = 0x000005A0; + result[8] = state1[0]; result[9] = state1[1]; result[10] = state1[2]; result[11] = state1[3]; + result[12] = state1[4]; result[13] = state1[5]; result[14] = state1[6]; result[15] = state1[7]; + sha256_round_body(in, result + 8); // inner length = 180 * 8 = 1184 = 0x5A0 + + in[0] = result[8]; in[1] = result[9]; in[2] = result[10]; in[3] = result[11]; + in[4] = result[12]; in[5] = result[13]; in[6] = result[14]; in[7] = result[15]; + in[8] = 0x80000000; in[15] = 0x00000300; + in[9] = in[10] = in[11] = in[12] = in[13] = in[14] = 0x00000000; + result[8] = state2[0]; result[9] = state2[1]; result[10] = state2[2]; result[11] = state2[3]; + result[12] = state2[4]; result[13] = state2[5]; result[14] = state2[6]; result[15] = state2[7]; + sha256_round_body(in, result + 8); // outer length = 96 * 8 = 768 = 0x300 + + *(uint4*)&Bdev(i, 0) = make_uint4(cuda_swab32(result[0]), cuda_swab32(result[5]), cuda_swab32(result[10]), cuda_swab32(result[15])); + *(uint4*)&Bdev(i, 4) = make_uint4(cuda_swab32(result[4]), cuda_swab32(result[9]), cuda_swab32(result[14]), cuda_swab32(result[3])); + *(uint4*)&Bdev(i, 8) = make_uint4(cuda_swab32(result[8]), cuda_swab32(result[13]), cuda_swab32(result[2]), cuda_swab32(result[7])); + *(uint4*)&Bdev(i, 12) = make_uint4(cuda_swab32(result[12]), cuda_swab32(result[1]), cuda_swab32(result[6]), cuda_swab32(result[11])); + + if (i == 0) { + sha256dev(0) = result[0]; + sha256dev(1) = result[1]; + sha256dev(2) = result[2]; + sha256dev(3) = result[3]; + sha256dev(4) = result[4]; + sha256dev(5) = result[5]; + sha256dev(6) = result[6]; + sha256dev(7) = result[7]; + } + } + } +} + __launch_bounds__(32, 1) __global__ void yescrypt_gpu_hash_k5(int threads, uint32_t startNonce, uint32_t *nonceVector, uint32_t target, const uint32_t r, const uint32_t p) { @@ -1153,7 +1248,7 @@ void yescrypt_cpu_init(int thr_id, int threads, uint32_t *d_hash1, uint32_t *d_h } __host__ -void yescrypt_setTarget(int thr_id, uint32_t pdata[20], char *key, uint32_t key_len) +void yescrypt_setTarget(int thr_id, uint32_t pdata[28], char *key, uint32_t key_len, const int perslen) { uint32_t h[8], data[32]; @@ -1170,7 +1265,7 @@ void yescrypt_setTarget(int thr_id, uint32_t pdata[20], char *key, uint32_t key_ sha256_round_body_host(data, h, cpu_K); cudaMemcpyToSymbol(cpu_h, h, 8 * sizeof(uint32_t), 0, cudaMemcpyHostToDevice); - cudaMemcpyToSymbol(c_data, pdata, 20 * sizeof(uint32_t), 0, cudaMemcpyHostToDevice); + cudaMemcpyToSymbol(c_data, pdata, 28 * sizeof(uint32_t), 0, cudaMemcpyHostToDevice); if (key) { @@ -1185,17 +1280,24 @@ void yescrypt_setTarget(int thr_id, uint32_t pdata[20], char *key, uint32_t key_ } else { - memcpy(data, pdata, 80); - data[20] = 0x80000000; - data[21] = data[22] = data[23] = data[24] = data[25] = data[26] = data[27] = data[28] = data[29] = data[30] = 0; - data[31] = (80 + 64) * 8; + if (perslen == 80) { + memcpy(data, pdata, 80); + data[20] = 0x80000000; + data[21] = data[22] = data[23] = data[24] = data[25] = data[26] = data[27] = data[28] = data[29] = data[30] = 0; + data[31] = (80 + 64) * 8; + } else { + memcpy(data, pdata, 112); + data[28] = 0x80000000; + data[29] = data[30] = 0; + data[31] = (112 + 64) * 8; + } key_len = 0; } CUDA_SAFE_CALL(cudaMemcpyToSymbol(client_key, data, 32 * sizeof(uint32_t), 0, cudaMemcpyHostToDevice)); CUDA_SAFE_CALL(cudaMemcpyToSymbol(client_key_len, &key_len, sizeof(uint32_t), 0, cudaMemcpyHostToDevice)); } -__host__ void yescrypt_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *resultnonces, uint32_t target, const uint32_t N, const uint32_t r, const uint32_t p) +__host__ void yescrypt_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *resultnonces, uint32_t target, const uint32_t N, const uint32_t r, const uint32_t p, const int is112) { int dev_id = device_map[thr_id % MAX_GPUS]; CUDA_SAFE_CALL(cudaMemset(d_GNonce[thr_id], 0, 2 * sizeof(uint32_t))); @@ -1236,7 +1338,10 @@ __host__ void yescrypt_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startN else if (device_sm[dev_id] > 300) loop_count = max(N * r / 4096, 1); else loop_count = max(N * r / 2048, 1); - yescrypt_gpu_hash_k0 << > > (threads, startNounce, r, p); + if (is112) + yescrypt_gpu_hash_k0_112bytes << > > (threads, startNounce, r, p); + else + yescrypt_gpu_hash_k0 << > > (threads, startNounce, r, p); CUDA_SAFE_CALL(cudaGetLastError()); for (uint32_t l = 0; l < p; l++) { diff --git a/yescrypt/yescrypt.cu b/yescrypt/yescrypt.cu index ddfcb14..e2966d7 100644 --- a/yescrypt/yescrypt.cu +++ b/yescrypt/yescrypt.cu @@ -10,8 +10,8 @@ extern "C" { } extern void yescrypt_cpu_init(int thr_id, int threads, uint32_t *d_hash1, uint32_t *d_hash2, uint32_t *d_hash3, uint32_t *d_hash4); -extern void yescrypt_setTarget(int thr_id, uint32_t pdata[20], char *key, uint32_t key_len); -extern void yescrypt_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *resultnonces, uint32_t target, const uint32_t N, const uint32_t r, const uint32_t p); +extern void yescrypt_setTarget(int thr_id, uint32_t pdata[20], char *key, uint32_t key_len, const int perslen); +extern void yescrypt_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *resultnonces, uint32_t target, const uint32_t N, const uint32_t r, const uint32_t p, const int is112); extern void yescrypt_cpu_free(int thr_id); extern char *yescrypt_key; @@ -20,50 +20,50 @@ extern uint32_t yescrypt_param_N; extern uint32_t yescrypt_param_r; extern uint32_t yescrypt_param_p; -void yescrypt_hash_base(void *state, const void *input, const uint32_t N, const uint32_t r, const uint32_t p, char *key, const size_t key_len) +void yescrypt_hash_base(void *state, const void *input, const uint32_t N, const uint32_t r, const uint32_t p, char *key, const size_t key_len, int perslen) { if (client_key_len == 0xff) { client_key = key; client_key_len = key_len; } - yescrypt_bsty((unsigned char*)input, 80, (unsigned char*)input, 80, N, r, p, (unsigned char *)state, 32); + yescrypt_bsty((unsigned char*)input, perslen, (unsigned char*)input, perslen, N, r, p, (unsigned char *)state, 32); } -void yescrypt_hash(void *state, const void *input) +void yescrypt_hash(void *state, const void *input, int perslen) { - yescrypt_hash_base(state, input, 2048, 8, 1, NULL, 0); + yescrypt_hash_base(state, input, 2048, 8, 1, NULL, 0, perslen); } void yescryptr8_hash(void *state, const void *input) { - yescrypt_hash_base(state, input, 2048, 8, 1, "Client Key", 10); + yescrypt_hash_base(state, input, 2048, 8, 1, (char *)"Client Key", 10, 80); } void yescryptr16_hash(void *state, const void *input) { - yescrypt_hash_base(state, input, 4096, 16, 1, "Client Key", 10); + yescrypt_hash_base(state, input, 4096, 16, 1, (char *)"Client Key", 10, 80); } void yescryptr16v2_hash(void *state, const void *input) { - yescrypt_hash_base(state, input, 4096, 16, 4, "PPTPPubKey", 10); + yescrypt_hash_base(state, input, 4096, 16, 4, (char *)"PPTPPubKey", 10, 80); } void yescryptr24_hash(void *state, const void *input) { - yescrypt_hash_base(state, input, 4096, 24, 1, "Jagaricoin", 10); + yescrypt_hash_base(state, input, 4096, 24, 1, (char *)"Jagaricoin", 10, 80); } void yescryptr32_hash(void *state, const void *input) { - yescrypt_hash_base(state, input, 4096, 32, 1, "WaviBanana", 10); + yescrypt_hash_base(state, input, 4096, 32, 1, (char *)"WaviBanana", 10, 80); } int scanhash_yescrypt_base(int thr_id, uint32_t *pdata, uint32_t *ptarget, uint32_t max_nonce, uint32_t *hashes_done, const uint32_t N, const uint32_t r, const uint32_t p, - char *key, const size_t key_len) { + char *key, const size_t key_len, int perslen) { static THREAD uint32_t *d_hash1 = nullptr; static THREAD uint32_t *d_hash2 = nullptr; static THREAD uint32_t *d_hash3 = nullptr; @@ -151,16 +151,16 @@ int scanhash_yescrypt_base(int thr_id, uint32_t *pdata, init = true; } - uint32_t endiandata[20]; - for (int k = 0; k < 20; k++) + uint32_t endiandata[32]; + for (int k = 0; k < 32; k++) be32enc(&endiandata[k], pdata[k]); - yescrypt_setTarget(thr_id, pdata, key, key_len); + yescrypt_setTarget(thr_id, pdata, key, key_len, perslen); do { uint32_t foundNonce[2] = { 0, 0 }; - yescrypt_cpu_hash_32(thr_id, throughput, pdata[19], foundNonce, ptarget[7], N, r, p); + yescrypt_cpu_hash_32(thr_id, throughput, pdata[19], foundNonce, ptarget[7], N, r, p, perslen == 112 ? 1 : 0); if (stop_mining) { @@ -176,7 +176,7 @@ int scanhash_yescrypt_base(int thr_id, uint32_t *pdata, if (opt_verify) { be32enc(&endiandata[19], foundNonce[0]); - yescrypt_hash_base(vhash64, endiandata, N, r, p, key, key_len); + yescrypt_hash_base(vhash64, endiandata, N, r, p, key, key_len, perslen); } if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) { int res = 1; @@ -186,11 +186,14 @@ int scanhash_yescrypt_base(int thr_id, uint32_t *pdata, if (opt_verify) { be32enc(&endiandata[19], foundNonce[1]); - yescrypt_hash_base(vhash64, endiandata, N, r, p, key, key_len); + yescrypt_hash_base(vhash64, endiandata, N, r, p, key, key_len, perslen); } if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) { - pdata[21] = foundNonce[1]; + if (perslen == 80) + pdata[21] = foundNonce[1]; + else + pdata[29] = foundNonce[1]; res++; if (opt_benchmark) applog(LOG_INFO, "GPU #%d Found second nonce %08x", thr_id, foundNonce[1]); } @@ -221,45 +224,45 @@ int scanhash_yescrypt_base(int thr_id, uint32_t *pdata, int scanhash_yescrypt(int thr_id, uint32_t *pdata, uint32_t *ptarget, uint32_t max_nonce, - uint32_t *hashes_done) + uint32_t *hashes_done, int perslen) { if (yescrypt_param_N == 0) yescrypt_param_N = 2048; if (yescrypt_param_r == 0) yescrypt_param_r = 8; if (yescrypt_param_p == 0) yescrypt_param_p = 1; - return scanhash_yescrypt_base(thr_id, pdata, ptarget, max_nonce, hashes_done, yescrypt_param_N, yescrypt_param_r, yescrypt_param_p, yescrypt_key, yescrypt_key_len); + return scanhash_yescrypt_base(thr_id, pdata, ptarget, max_nonce, hashes_done, yescrypt_param_N, yescrypt_param_r, yescrypt_param_p, yescrypt_key, yescrypt_key_len, perslen); } int scanhash_yescryptr8(int thr_id, uint32_t *pdata, uint32_t *ptarget, uint32_t max_nonce, uint32_t *hashes_done) { - return scanhash_yescrypt_base(thr_id, pdata, ptarget, max_nonce, hashes_done, 2048, 8, 1, "Client Key", 10); + return scanhash_yescrypt_base(thr_id, pdata, ptarget, max_nonce, hashes_done, 2048, 8, 1, (char *)"Client Key", 10, 80); } int scanhash_yescryptr16(int thr_id, uint32_t *pdata, uint32_t *ptarget, uint32_t max_nonce, uint32_t *hashes_done) { - return scanhash_yescrypt_base(thr_id, pdata, ptarget, max_nonce, hashes_done, 4096, 16, 1, "Client Key", 10); + return scanhash_yescrypt_base(thr_id, pdata, ptarget, max_nonce, hashes_done, 4096, 16, 1, (char *)"Client Key", 10, 80); } int scanhash_yescryptr16v2(int thr_id, uint32_t *pdata, uint32_t *ptarget, uint32_t max_nonce, uint32_t *hashes_done) { - return scanhash_yescrypt_base(thr_id, pdata, ptarget, max_nonce, hashes_done, 4096, 16, 4, "PPTPPubKey", 10); + return scanhash_yescrypt_base(thr_id, pdata, ptarget, max_nonce, hashes_done, 4096, 16, 4, (char *)"PPTPPubKey", 10, 80); } int scanhash_yescryptr24(int thr_id, uint32_t *pdata, uint32_t *ptarget, uint32_t max_nonce, uint32_t *hashes_done) { - return scanhash_yescrypt_base(thr_id, pdata, ptarget, max_nonce, hashes_done, 4096, 24, 1, "Jagaricoin", 10); + return scanhash_yescrypt_base(thr_id, pdata, ptarget, max_nonce, hashes_done, 4096, 24, 1, (char *)"Jagaricoin", 10, 80); } int scanhash_yescryptr32(int thr_id, uint32_t *pdata, uint32_t *ptarget, uint32_t max_nonce, uint32_t *hashes_done) { - return scanhash_yescrypt_base(thr_id, pdata, ptarget, max_nonce, hashes_done, 4096, 32, 1, "WaviBanana", 10); + return scanhash_yescrypt_base(thr_id, pdata, ptarget, max_nonce, hashes_done, 4096, 32, 1, (char *)"WaviBanana", 10, 80); }