fix high cpu usage for mtp-tcr solo and mtp-classic

compat/ccminer-config.h

@@ -164,7 +164,7 @@
 #define PACKAGE_URL "http://github.com/zcoinofficial/ccminer"
 
 /* Define to the version of this package. */
-#define PACKAGE_VERSION "1.2.9-djm34-beta"
+#define PACKAGE_VERSION "1.2.9-djm34"
 
 /* If using the C implementation of alloca, define if you know the
    direction of stack growth for your system; otherwise it will be

configure.ac

@@ -1,4 +1,4 @@
-AC_INIT([ccminer], [1.2.9L-djm34-beta], [], [ccminer], [http://github.com/zcoinofficial/ccminer])
+AC_INIT([ccminer], [1.2.9L-djm34], [], [ccminer], [http://github.com/zcoinofficial/ccminer])
 
 AC_PREREQ([2.59c])
 AC_CANONICAL_SYSTEM

cuda_mtp/cuda_mtp_forlib.cu

@@ -986,7 +986,7 @@ uint32_t get_tpb_mtp(int thr_id)
 
 
 __host__
-void mtp_setBlockTarget(int thr_id, const void* pDataIn, const void *pTargetIn, const void * zElement)
+void mtp_setBlockTarget_old(int thr_id, const void* pDataIn, const void *pTargetIn, const void * zElement)
 {
 //	cudaSetDevice(device_map[thr_id]);
 
@@ -997,7 +997,7 @@ void mtp_setBlockTarget(int thr_id, const void* pDataIn, const void *pTargetIn,
 }
 
 __host__
-void mtp_setBlockTarget_test(int thr_id, const void* pDataIn, const void *pTargetIn, const void * zElement,cudaStream_t s0)
+void mtp_setBlockTarget(int thr_id, const void* pDataIn, const void *pTargetIn, const void * zElement,cudaStream_t s0)
 {
 	//	cudaSetDevice(device_map[thr_id]);
 
@@ -1024,12 +1024,11 @@ void mtp_fill(uint32_t dev_id, const uint64_t *Block, uint32_t offset, uint32_t
 }
 
 __host__
-uint32_t mtp_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startNounce)
+uint32_t mtp_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startNounce,cudaStream_t s0)
 {
 //	cudaSetDevice(device_map[thr_id]);
 	uint32_t result = UINT32_MAX;
-	CUDA_SAFE_CALL(cudaMemset(d_MinNonces[thr_id], 0xff, sizeof(uint32_t)));
-//	int dev_id = device_map[thr_id % MAX_GPUS];
+	cudaMemsetAsync(d_MinNonces[thr_id], 0xff, sizeof(uint32_t), s0);
 
 
 	uint32_t tpb = TPB_MTP; //TPB52;
@@ -1040,11 +1039,10 @@ uint32_t mtp_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startNounce)
 	dim3 blockyloop(tpb);
 
 	//yloop_init <<<gridyloop, blockyloop>>>(thr_id, threads, startNounce, GYLocal[thr_id]);
+	cudaStreamSynchronize(s0);
 
 	mtp_yloop << < gridyloop, blockyloop >> >(thr_id, threads, startNounce, (Type*)HBlock[thr_id],  d_MinNonces[thr_id]);
-
-
-	CUDA_SAFE_CALL(cudaMemcpy(h_MinNonces[thr_id], d_MinNonces[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost));
+	cudaStreamSynchronize(s0);
 
 	result = *h_MinNonces[thr_id];
 	return result; 
@@ -1785,11 +1783,11 @@ __global__  void mtp_fc2(uint32_t threads, uint4  *  DBlock, uint2 *a) {
 
 
 
-__host__ void get_tree(int thr_id, uint8_t* d) {
+__host__ void get_tree_old(int thr_id, uint8_t* d) {
 	CUDA_SAFE_CALL(cudaMemcpy(d, buffer_a[thr_id], sizeof(uint2) * 2 * 1048576 * 4, cudaMemcpyDeviceToHost));
 }
 
-__host__ void get_tree_test(int thr_id, uint8_t* d, cudaStream_t s0) {
+__host__ void get_tree(int thr_id, uint8_t* d, cudaStream_t s0) {
 	CUDA_SAFE_CALL(cudaMemcpyAsync(d, buffer_a[thr_id], sizeof(uint2) * 2 * 1048576 * 4, cudaMemcpyDeviceToHost, s0));
 }
 
@@ -1913,7 +1911,7 @@ void mtp_fill_1b(int thr_id, uint64_t *Block, uint32_t block_nr)
 }
 
 __host__
-void mtp_fill_1c_test(int thr_id, uint64_t *Block, uint32_t block_nr, cudaStream_t s0)
+void mtp_fill_1c(int thr_id, uint64_t *Block, uint32_t block_nr, cudaStream_t s0)
 {
 //	cudaSetDevice(device_map[thr_id]);
 	//	uint4 *Blockptr = &HBlock[thr_id][block_nr * 64];
@@ -1934,7 +1932,7 @@ void mtp_fill_1c_test(int thr_id, uint64_t *Block, uint32_t block_nr, cudaStream
 }
 
 __host__
-void mtp_fill_1c(int thr_id, uint64_t *Block, uint32_t block_nr)
+void mtp_fill_1c_old(int thr_id, uint64_t *Block, uint32_t block_nr)
 {
 	//	cudaSetDevice(device_map[thr_id]);
 	//	uint4 *Blockptr = &HBlock[thr_id][block_nr * 64];

cuda_mtp/mtp-tcr.cu

@@ -13,22 +13,13 @@
 #define memcost 4*1024*1024
 
 extern void mtp_cpu_init(int thr_id, uint32_t threads);
-
 extern uint32_t mtptcr_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startNounce, cudaStream_t s0);
-
-extern void mtp_setBlockTarget(int thr_id,const void* pDataIn, const void *pTargetIn, const void * zElement);
-extern void mtp_setBlockTarget_test(int thr_id, const void* pDataIn, const void *pTargetIn, const void * zElement,cudaStream_t s0);
-extern void mtp_fill(uint32_t d, const uint64_t *Block, uint32_t offset, uint32_t datachunk);
-extern void mtp_fill_1b(int thr_id, uint64_t *Block, uint32_t block_nr);
+extern void mtp_setBlockTarget(int thr_id, const void* pDataIn, const void *pTargetIn, const void * zElement,cudaStream_t s0);
 extern uint32_t get_tpb_mtp(int thr_id);
-
-extern void mtp_fill_1c(int thr_id, uint64_t *Block, uint32_t block_nr);
-extern void mtp_fill_1c_test(int thr_id, uint64_t *Block, uint32_t block_nr, cudaStream_t s0);
-//extern void mtp_i_cpu(int thr_id, uint32_t *block_header);
-extern void mtp_i_cpu(int thr_id, uint32_t *block_header, cudaStream_t s0);
+extern void mtp_fill_1c(int thr_id, uint64_t *Block, uint32_t block_nr, cudaStream_t s0);
 extern void mtp_i_cpu2(int thr_id, uint32_t *block_header, cudaStream_t s0);
-void get_tree(int thr_id, uint8_t* d);
-void get_tree_test(int thr_id, uint8_t* d, cudaStream_t s0);
+void get_tree(int thr_id, uint8_t* d, cudaStream_t s0);
+
 #define HASHLEN 32
 #define SALTLEN 16
 #define PWD "password"
@@ -60,7 +51,7 @@ extern "C" int scanhash_mtptcr(int nthreads,int thr_id, struct work* work, uint3
 //if (JobId==0)
 //	pthread_barrier_init(&barrier, NULL, nthreads);
 
-		cudaStream_t s0;
+	cudaStream_t s0;
 	uint32_t *pdata = work->data;
 	uint32_t *ptarget = work->target;
 
@@ -131,18 +122,18 @@ if (JobId[thr_id] != work->data[16] || XtraNonce2[thr_id] != ((uint64_t*)work->x
 context[thr_id] = init_argon2d_param((const char*)endiandata);
 
 argon2_ctx_from_mtp(&context[thr_id], &instance[thr_id]);
-	mtp_fill_1c_test(thr_id, instance[thr_id].memory[0 + 0].v, 0 + 0,s0);
-	mtp_fill_1c_test(thr_id, instance[thr_id].memory[0 + 1].v, 0 + 1,s0);
-	mtp_fill_1c_test(thr_id, instance[thr_id].memory[2 + 0].v, 1048576 + 0,s0);
-	mtp_fill_1c_test(thr_id, instance[thr_id].memory[2 + 1].v, 1048576 + 1,s0);
-	mtp_fill_1c_test(thr_id, instance[thr_id].memory[4 + 0].v, 2097152 + 0,s0);
-	mtp_fill_1c_test(thr_id, instance[thr_id].memory[4 + 1].v, 2097152 + 1,s0);
-	mtp_fill_1c_test(thr_id, instance[thr_id].memory[6 + 0].v, 3145728 + 0,s0);
-	mtp_fill_1c_test(thr_id, instance[thr_id].memory[6 + 1].v, 3145728 + 1,s0);
+	mtp_fill_1c(thr_id, instance[thr_id].memory[0 + 0].v, 0 + 0,s0);
+	mtp_fill_1c(thr_id, instance[thr_id].memory[0 + 1].v, 0 + 1,s0);
+	mtp_fill_1c(thr_id, instance[thr_id].memory[2 + 0].v, 1048576 + 0,s0);
+	mtp_fill_1c(thr_id, instance[thr_id].memory[2 + 1].v, 1048576 + 1,s0);
+	mtp_fill_1c(thr_id, instance[thr_id].memory[4 + 0].v, 2097152 + 0,s0);
+	mtp_fill_1c(thr_id, instance[thr_id].memory[4 + 1].v, 2097152 + 1,s0);
+	mtp_fill_1c(thr_id, instance[thr_id].memory[6 + 0].v, 3145728 + 0,s0);
+	mtp_fill_1c(thr_id, instance[thr_id].memory[6 + 1].v, 3145728 + 1,s0);
 
 	mtp_i_cpu2(thr_id, instance[thr_id].block_header,s0);
 
-	get_tree_test(thr_id,dx[thr_id],s0);
+	get_tree(thr_id,dx[thr_id],s0);
 
 	cudaStreamSynchronize(s0);
 
@@ -154,7 +145,7 @@ argon2_ctx_from_mtp(&context[thr_id], &instance[thr_id]);
 
 	std::copy(root.begin(), root.end(), TheMerkleRoot[thr_id]);
 
-	mtp_setBlockTarget_test(thr_id, endiandata, ptarget, &TheMerkleRoot[thr_id],s0);
+	mtp_setBlockTarget(thr_id, endiandata, ptarget, &TheMerkleRoot[thr_id],s0);
 
 	root.resize(0);
 }
@@ -185,7 +176,7 @@ argon2_ctx_from_mtp(&context[thr_id], &instance[thr_id]);
 			blockS nBlockMTP[MTP_L *2] = {0};
 			unsigned char nProofMTP[MTP_L * 3 * 353 ] = {0};
 
-			uint32_t is_sol = mtptcr_solver_test(thr_id,foundNonce, &instance[thr_id], nBlockMTP,nProofMTP, TheMerkleRoot[thr_id], mtpHashValue, *ordered_tree[thr_id], endiandata,TheUint256Target[0],s0);
+			uint32_t is_sol = mtptcr_solver(thr_id,foundNonce, &instance[thr_id], nBlockMTP,nProofMTP, TheMerkleRoot[thr_id], mtpHashValue, *ordered_tree[thr_id], endiandata,TheUint256Target[0],s0);
 
 			if (JobId[thr_id] != work->data[16] || XtraNonce2[thr_id] != ((uint64_t*)work->xnonce2)[0])
 				return 0; // if work has changed stop and go back to the initialization
@@ -318,19 +309,20 @@ extern "C" int scanhash_mtptcr_solo(int nthreads, int thr_id, struct work* work,
 
 		argon2_ctx_from_mtp(&context[thr_id], &instance[thr_id]);
 
-		mtp_fill_1c(thr_id, instance[thr_id].memory[0 + 0].v, 0 + 0);
-		mtp_fill_1c(thr_id, instance[thr_id].memory[0 + 1].v, 0 + 1);
-		mtp_fill_1c(thr_id, instance[thr_id].memory[2 + 0].v, 1048576 + 0);
-		mtp_fill_1c(thr_id, instance[thr_id].memory[2 + 1].v, 1048576 + 1);
-		mtp_fill_1c(thr_id, instance[thr_id].memory[4 + 0].v, 2097152 + 0);
-		mtp_fill_1c(thr_id, instance[thr_id].memory[4 + 1].v, 2097152 + 1);
-		mtp_fill_1c(thr_id, instance[thr_id].memory[6 + 0].v, 3145728 + 0);
-		mtp_fill_1c(thr_id, instance[thr_id].memory[6 + 1].v, 3145728 + 1);
+		mtp_fill_1c(thr_id, instance[thr_id].memory[0 + 0].v, 0 + 0,s0);
+		mtp_fill_1c(thr_id, instance[thr_id].memory[0 + 1].v, 0 + 1,s0);
+		mtp_fill_1c(thr_id, instance[thr_id].memory[2 + 0].v, 1048576 + 0,s0);
+		mtp_fill_1c(thr_id, instance[thr_id].memory[2 + 1].v, 1048576 + 1,s0);
+		mtp_fill_1c(thr_id, instance[thr_id].memory[4 + 0].v, 2097152 + 0,s0);
+		mtp_fill_1c(thr_id, instance[thr_id].memory[4 + 1].v, 2097152 + 1,s0);
+		mtp_fill_1c(thr_id, instance[thr_id].memory[6 + 0].v, 3145728 + 0,s0);
+		mtp_fill_1c(thr_id, instance[thr_id].memory[6 + 1].v, 3145728 + 1,s0);
 
 		mtp_i_cpu2(thr_id, instance[thr_id].block_header,s0);
 
-		get_tree(thr_id, dx[thr_id]);
+		get_tree(thr_id, dx[thr_id],s0);
 
+		cudaStreamSynchronize(s0);
 		//	printf("Step 2 : Compute the root Φ of the Merkle hash tree \n");
 		//  sleep(10);
 
@@ -342,7 +334,7 @@ extern "C" int scanhash_mtptcr_solo(int nthreads, int thr_id, struct work* work,
 
 		std::copy(root.begin(), root.end(), TheMerkleRoot[thr_id]);
 
-		mtp_setBlockTarget(thr_id, endiandata, ptarget, &TheMerkleRoot[thr_id]);
+		mtp_setBlockTarget(thr_id, endiandata, ptarget, &TheMerkleRoot[thr_id],s0);
 		root.resize(0);
 	}
 
@@ -371,7 +363,7 @@ extern "C" int scanhash_mtptcr_solo(int nthreads, int thr_id, struct work* work,
 			blockS nBlockMTP[MTP_L * 2] = { 0 };
 			unsigned char nProofMTP[MTP_L * 3 * 353] = { 0 };
 
-			uint32_t is_sol = mtptcr_solver(thr_id, foundNonce, &instance[thr_id], nBlockMTP, nProofMTP, TheMerkleRoot[thr_id], mtpHashValue, *ordered_tree[thr_id], endiandata, TheUint256Target[0]);
+			uint32_t is_sol = mtptcr_solver(thr_id, foundNonce, &instance[thr_id], nBlockMTP, nProofMTP, TheMerkleRoot[thr_id], mtpHashValue, *ordered_tree[thr_id], endiandata, TheUint256Target[0],s0);
 
 			if (is_sol == 1 /*&& fulltest(vhash64, ptarget)*/) {
 

cuda_mtp/mtp.cu

@@ -10,20 +10,13 @@
 #define memcost 4*1024*1024
 
 extern void mtp_cpu_init(int thr_id, uint32_t threads);
-
-extern uint32_t mtp_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startNounce);
-
-extern void mtp_setBlockTarget(int thr_id,const void* pDataIn, const void *pTargetIn, const void * zElement);
-extern void mtp_fill(uint32_t d, const uint64_t *Block, uint32_t offset, uint32_t datachunk);
-extern void mtp_fill_1b(int thr_id, uint64_t *Block, uint32_t block_nr);
+extern uint32_t mtp_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startNounce, cudaStream_t s0);
+extern void mtp_setBlockTarget(int thr_id,const void* pDataIn, const void *pTargetIn, const void * zElement, cudaStream_t s0);
 extern uint32_t get_tpb_mtp(int thr_id);
-
-extern void mtp_fill_1c(int thr_id, uint64_t *Block, uint32_t block_nr);
-
-//extern void mtp_i_cpu(int thr_id, uint32_t *block_header);
-extern void mtp_i_cpu(int thr_id, uint32_t *block_header, cudaStream_t s0);
+extern void mtp_fill_1c(int thr_id, uint64_t *Block, uint32_t block_nr, cudaStream_t s0);
 extern void mtp_i_cpu2(int thr_id, uint32_t *block_header, cudaStream_t s0);
-void get_tree(int thr_id, uint8_t* d);
+void get_tree(int thr_id, uint8_t* d, cudaStream_t s0);
+
 #define HASHLEN 32
 #define SALTLEN 16
 #define PWD "password" 
@@ -51,11 +44,11 @@ extern "C" int scanhash_mtp(int nthreads,int thr_id, struct work* work, uint32_t
 {
 
 	unsigned char mtpHashValue[32];
-	cudaStream_t s0;
+
 //if (JobId==0)
 //	pthread_barrier_init(&barrier, NULL, nthreads);
 
-
+	cudaStream_t s0;
 	uint32_t *pdata = work->data;
 	uint32_t *ptarget = work->target;
 
@@ -124,22 +117,20 @@ if (JobId[thr_id] != work->data[16] || XtraNonce2[thr_id] != ((uint64_t*)work->x
 context[thr_id] = init_argon2d_param((const char*)endiandata);
 
 argon2_ctx_from_mtp(&context[thr_id], &instance[thr_id]);
-	mtp_fill_1c(thr_id, instance[thr_id].memory[0 + 0].v, 0 + 0);
-	mtp_fill_1c(thr_id, instance[thr_id].memory[0 + 1].v, 0 + 1);
-	mtp_fill_1c(thr_id, instance[thr_id].memory[2 + 0].v, 1048576 + 0);
-	mtp_fill_1c(thr_id, instance[thr_id].memory[2 + 1].v, 1048576 + 1);
-	mtp_fill_1c(thr_id, instance[thr_id].memory[4 + 0].v, 2097152 + 0);
-	mtp_fill_1c(thr_id, instance[thr_id].memory[4 + 1].v, 2097152 + 1);
-	mtp_fill_1c(thr_id, instance[thr_id].memory[6 + 0].v, 3145728 + 0);
-	mtp_fill_1c(thr_id, instance[thr_id].memory[6 + 1].v, 3145728 + 1);
+	mtp_fill_1c(thr_id, instance[thr_id].memory[0 + 0].v, 0 + 0,s0);
+	mtp_fill_1c(thr_id, instance[thr_id].memory[0 + 1].v, 0 + 1,s0);
+	mtp_fill_1c(thr_id, instance[thr_id].memory[2 + 0].v, 1048576 + 0,s0);
+	mtp_fill_1c(thr_id, instance[thr_id].memory[2 + 1].v, 1048576 + 1,s0);
+	mtp_fill_1c(thr_id, instance[thr_id].memory[4 + 0].v, 2097152 + 0,s0);
+	mtp_fill_1c(thr_id, instance[thr_id].memory[4 + 1].v, 2097152 + 1,s0);
+	mtp_fill_1c(thr_id, instance[thr_id].memory[6 + 0].v, 3145728 + 0,s0);
+	mtp_fill_1c(thr_id, instance[thr_id].memory[6 + 1].v, 3145728 + 1,s0);
 
 	mtp_i_cpu2(thr_id, instance[thr_id].block_header,s0);
 
+	get_tree(thr_id,dx[thr_id],s0);
 
-
-	get_tree(thr_id,dx[thr_id]);
-
-
+	cudaStreamSynchronize(s0);
 //  sleep(10);
 
 	ordered_tree[thr_id] = new MerkleTree(dx[thr_id], true);
@@ -150,7 +141,7 @@ argon2_ctx_from_mtp(&context[thr_id], &instance[thr_id]);
 
 	std::copy(root.begin(), root.end(), TheMerkleRoot[thr_id]);
 
-	mtp_setBlockTarget(thr_id, endiandata, ptarget, &TheMerkleRoot[thr_id]);
+	mtp_setBlockTarget(thr_id, endiandata, ptarget, &TheMerkleRoot[thr_id],s0);
 	root.resize(0);
 }
 
@@ -162,7 +153,7 @@ argon2_ctx_from_mtp(&context[thr_id], &instance[thr_id]);
 //		cudaProfilerStart();
 //		cudaProfilerStop();
 		*hashes_done = pdata[19] - first_nonce + throughput;
-		foundNonce = mtp_cpu_hash_32(thr_id, throughput, pdata[19]);
+		foundNonce = mtp_cpu_hash_32(thr_id, throughput, pdata[19],s0);
 //		cudaProfilerStop();
 		uint32_t _ALIGN(64) vhash64[8];
 		if (foundNonce != UINT32_MAX)
@@ -181,7 +172,7 @@ argon2_ctx_from_mtp(&context[thr_id], &instance[thr_id]);
 			blockS nBlockMTP[MTP_L *2] = {0};
 			unsigned char nProofMTP[MTP_L * 3 * 353 ] = {0};
 
-			uint32_t is_sol = mtp_solver(thr_id,foundNonce, &instance[thr_id], nBlockMTP,nProofMTP, TheMerkleRoot[thr_id], mtpHashValue, *ordered_tree[thr_id], endiandata,TheUint256Target[0]);
+			uint32_t is_sol = mtp_solver(thr_id,foundNonce, &instance[thr_id], nBlockMTP,nProofMTP, TheMerkleRoot[thr_id], mtpHashValue, *ordered_tree[thr_id], endiandata,TheUint256Target[0],s0);
 
 			if (is_sol==1 /*&& fulltest(vhash64, ptarget)*/) {
 
@@ -315,22 +306,22 @@ extern "C" int scanhash_mtp_solo(int nthreads, int thr_id, struct work* work, ui
 
 		argon2_ctx_from_mtp(&context[thr_id], &instance[thr_id]);
 
-		mtp_fill_1c(thr_id, instance[thr_id].memory[0 + 0].v, 0 + 0);
-		mtp_fill_1c(thr_id, instance[thr_id].memory[0 + 1].v, 0 + 1);
-		mtp_fill_1c(thr_id, instance[thr_id].memory[2 + 0].v, 1048576 + 0);
-		mtp_fill_1c(thr_id, instance[thr_id].memory[2 + 1].v, 1048576 + 1);
-		mtp_fill_1c(thr_id, instance[thr_id].memory[4 + 0].v, 2097152 + 0);
-		mtp_fill_1c(thr_id, instance[thr_id].memory[4 + 1].v, 2097152 + 1);
-		mtp_fill_1c(thr_id, instance[thr_id].memory[6 + 0].v, 3145728 + 0);
-		mtp_fill_1c(thr_id, instance[thr_id].memory[6 + 1].v, 3145728 + 1);
+		mtp_fill_1c(thr_id, instance[thr_id].memory[0 + 0].v, 0 + 0,s0);
+		mtp_fill_1c(thr_id, instance[thr_id].memory[0 + 1].v, 0 + 1,s0);
+		mtp_fill_1c(thr_id, instance[thr_id].memory[2 + 0].v, 1048576 + 0,s0);
+		mtp_fill_1c(thr_id, instance[thr_id].memory[2 + 1].v, 1048576 + 1,s0);
+		mtp_fill_1c(thr_id, instance[thr_id].memory[4 + 0].v, 2097152 + 0,s0);
+		mtp_fill_1c(thr_id, instance[thr_id].memory[4 + 1].v, 2097152 + 1,s0);
+		mtp_fill_1c(thr_id, instance[thr_id].memory[6 + 0].v, 3145728 + 0,s0);
+		mtp_fill_1c(thr_id, instance[thr_id].memory[6 + 1].v, 3145728 + 1,s0);
 
 		mtp_i_cpu2(thr_id, instance[thr_id].block_header,s0);
 
-		get_tree(thr_id, dx[thr_id]);
+		get_tree(thr_id, dx[thr_id],s0);
 
 		//	printf("Step 2 : Compute the root Φ of the Merkle hash tree \n");
 		//  sleep(10);
-
+		cudaStreamSynchronize(s0);
 		ordered_tree[thr_id] = new MerkleTree(dx[thr_id], true);
 
 		JobId[thr_id] = work->data[17];
@@ -339,7 +330,7 @@ extern "C" int scanhash_mtp_solo(int nthreads, int thr_id, struct work* work, ui
 
 		std::copy(root.begin(), root.end(), TheMerkleRoot[thr_id]);
 
-		mtp_setBlockTarget(thr_id, endiandata, ptarget, &TheMerkleRoot[thr_id]);
+		mtp_setBlockTarget(thr_id, endiandata, ptarget, &TheMerkleRoot[thr_id],s0);
 		root.resize(0);
 	}
 
@@ -355,7 +346,7 @@ extern "C" int scanhash_mtp_solo(int nthreads, int thr_id, struct work* work, ui
 		uint32_t foundNonce;
 
 		*hashes_done = pdata[19] - first_nonce + throughput;
-		foundNonce = mtp_cpu_hash_32(thr_id, throughput, pdata[19]);
+		foundNonce = mtp_cpu_hash_32(thr_id, throughput, pdata[19],s0);
 
 		uint32_t _ALIGN(64) vhash64[8];
 		if (foundNonce != UINT32_MAX)
@@ -368,7 +359,7 @@ extern "C" int scanhash_mtp_solo(int nthreads, int thr_id, struct work* work, ui
 			blockS nBlockMTP[MTP_L * 2] = { 0 };
 			unsigned char nProofMTP[MTP_L * 3 * 353] = { 0 };
 
-			uint32_t is_sol = mtp_solver(thr_id, foundNonce, &instance[thr_id], nBlockMTP, nProofMTP, TheMerkleRoot[thr_id], mtpHashValue, *ordered_tree[thr_id], endiandata, TheUint256Target[0]);
+			uint32_t is_sol = mtp_solver(thr_id, foundNonce, &instance[thr_id], nBlockMTP, nProofMTP, TheMerkleRoot[thr_id], mtpHashValue, *ordered_tree[thr_id], endiandata, TheUint256Target[0],s0);
 
 			if (is_sol == 1 /*&& fulltest(vhash64, ptarget)*/) {