Skip to content

Instantly share code, notes, and snippets.

@wo01

wo01/sapling.diff

Last active January 1, 2019 02:51
Show Gist options
  • Star 0 You must be signed in to star a gist
  • Fork 0 You must be signed in to fork a gist
  • Save wo01/75ecb9660f5848ab308f93ce6cd42921 to your computer and use it in GitHub Desktop.
Save wo01/75ecb9660f5848ab308f93ce6cd42921 to your computer and use it in GitHub Desktop.
Download ZIP
ccminer-KlausT-8.21-mod-r17 patch for Koto Sapling
Raw
sapling.diff
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];
@@ -503,6 +503,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)
{
int thread = (blockDim.x * blockIdx.x + threadIdx.x);
@@ -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 << <grid, block >> > (threads, startNounce, r, p);
+ if (is112)
+ yescrypt_gpu_hash_k0_112bytes << <grid, block >> > (threads, startNounce, r, p);
+ else
+ yescrypt_gpu_hash_k0 << <grid, block >> > (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);
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment