faster kernels choice between --kernel yescrypt and --kernel yescrypt-multi

--kernel yescrypt for amd
--kernel yescrypt-multi for nvidia (> 5.0)

algorithm.c

@@ -58,7 +58,8 @@ const char *algorithm_type_str[] = {
   "Neoscrypt",
   "Lyra2RE",
   "pluck",
-  "yescrypt"
+  "yescrypt",
+  "yescrypt-multi"
 };
 
 void sha256(const unsigned char *message, unsigned int len, unsigned char *digest)
@@ -196,6 +197,8 @@ static cl_int queue_pluck_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_un
 	le_target = (cl_uint)le32toh(((uint32_t *)blk->work->/*device_*/target)[7]);
 //	memcpy(clState->cldata, blk->work->data, 80);
 	flip80(clState->cldata, blk->work->data);
+//int i;
+//for (i = 0; i<20; i++) ((uint32_t*)clState->cldata)[i] = ((uint32_t*)blk->work->data)[i]; // don't flip
 	status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 80, clState->cldata, 0, NULL, NULL);
 
 	CL_SET_ARG(clState->CLbuffer0);
@@ -215,9 +218,12 @@ static cl_int queue_yescrypt_kernel(_clState *clState, dev_blk_ctx *blk, __maybe
 	cl_int status = 0;
 
 
-	//	le_target = (*(cl_uint *)(blk->work->device_target + 28));
+//	le_target = (*(cl_uint *)(blk->work->device_target + 28));
 	le_target = (cl_uint)le32toh(((uint32_t *)blk->work->/*device_*/target)[7]);
-	//	memcpy(clState->cldata, blk->work->data, 80);
+//	le_target = (cl_uint)((uint32_t *)blk->work->target)[7];
+
+
+//	memcpy(clState->cldata, blk->work->data, 80);
 	flip80(clState->cldata, blk->work->data);
 	status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 80, clState->cldata, 0, NULL, NULL);
 
@@ -225,6 +231,7 @@ static cl_int queue_yescrypt_kernel(_clState *clState, dev_blk_ctx *blk, __maybe
 	CL_SET_ARG(clState->outputBuffer);
 	CL_SET_ARG(clState->padbuffer8);
 	CL_SET_ARG(clState->buffer1);
+	CL_SET_ARG(clState->buffer2);
 	CL_SET_ARG(le_target);
 
 	return status;
@@ -241,29 +248,54 @@ static cl_int queue_yescrypt_multikernel(_clState *clState, dev_blk_ctx *blk, __
 
 	//	le_target = (*(cl_uint *)(blk->work->device_target + 28));
 	le_target = (cl_uint)le32toh(((uint32_t *)blk->work->/*device_*/target)[7]);
-	//	memcpy(clState->cldata, blk->work->data, 80);
+	memcpy(clState->cldata, blk->work->data, 80);
-	flip80(clState->cldata, blk->work->data);
+//	flip80(clState->cldata, blk->work->data);
 	status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 80, clState->cldata, 0, NULL, NULL);
 //pbkdf and initial sha
 	kernel = &clState->kernel;
-	CL_SET_ARG_0(clState->CLbuffer0);
+
-	CL_SET_ARG(clState->buffer2);
+	CL_SET_ARG(clState->CLbuffer0);
-	CL_SET_ARG(clState->buffer3);
+	CL_SET_ARG(clState->outputBuffer);
-//mix1_1 (salsa)
+	CL_SET_ARG(clState->padbuffer8);
-	num = 0;
-	kernel = clState->extra_kernels;
-	CL_SET_ARG_0(clState->buffer1);
-	CL_SET_ARG(clState->buffer2);
-//mix1_2/2_2 (pwxform)
-	CL_NEXTKERNEL_SET_ARG_0(clState->padbuffer8);
 	CL_SET_ARG(clState->buffer1);
 	CL_SET_ARG(clState->buffer2);
+	CL_SET_ARG(clState->buffer3);
+	CL_SET_ARG(le_target);
+
+//inactive kernel
+	num = 0;
+	kernel = clState->extra_kernels;
+	CL_SET_ARG_N(0,clState->buffer1);
+	CL_SET_ARG_N(1,clState->buffer2);
+//	CL_SET_ARG_N(3, clState->buffer3);
+
 //mix2_2
-//	CL_NEXTKERNEL_SET_ARG_0(clState->padbuffer8);
+	num = 0;
-//	CL_SET_ARG(clState->buffer1);
+	CL_NEXTKERNEL_SET_ARG_N(0, clState->padbuffer8);
-//	CL_SET_ARG(clState->buffer2);
+	CL_SET_ARG_N(1,clState->buffer1);
+	CL_SET_ARG_N(2,clState->buffer2);
+	//mix2_2
+//inactive kernel
+//	num = 0;
+//	CL_NEXTKERNEL_SET_ARG_N(0, clState->buffer1);
+//	CL_SET_ARG_N(1, clState->buffer2);
+	//mix2_2
+
+	num = 0;
+	CL_NEXTKERNEL_SET_ARG_N(0, clState->padbuffer8);
+	CL_SET_ARG_N(1, clState->buffer1);
+	CL_SET_ARG_N(2, clState->buffer2);
+
+	//inactive kernel
+//	num = 0;
+//	CL_NEXTKERNEL_SET_ARG_N(0, clState->buffer1);
+//	CL_SET_ARG_N(1, clState->buffer2);
+	//mix2_2
+
+
 //pbkdf and finalization
-	CL_NEXTKERNEL_SET_ARG_0(clState->CLbuffer0);
+    num=0;
+	CL_NEXTKERNEL_SET_ARG(clState->CLbuffer0);
 	CL_SET_ARG(clState->outputBuffer);
 	CL_SET_ARG(clState->buffer2);
 	CL_SET_ARG(clState->buffer3);
@@ -819,10 +851,10 @@ static algorithm_settings_t algos[] = {
   A_YESCRYPT("yescrypt"),
 #undef A_YESCRYPT
 
-//#define A_YESCRYPT(a) \
+#define A_YESCRYPT_MULTI(a) \
-//                  { a, ALGO_YESCRYPT, "", 1, 65536, 65536, 0, 0, 0xFF, 0xFFFF000000000000ULL, 0x0000ffffUL, 3, -1,CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE , yescrypt_regenhash, queue_yescrypt_multikernel, gen_hash, append_neoscrypt_compiler_options}
+                  { a, ALGO_YESCRYPT_MULTI, "", 1, 65536, 65536, 0, 0, 0xFF, 0xFFFF000000000000ULL, 0x0000ffffUL, 4,-1,CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE , yescrypt_regenhash, queue_yescrypt_multikernel, gen_hash, append_neoscrypt_compiler_options}
-//  A_YESCRYPT("yescrypt"),
+  A_YESCRYPT_MULTI("yescrypt-multi"),
-//#undef A_YESCRYPT
+#undef A_YESCRYPT_MULTI
 
 
   // kernels starting from this will have difficulty calculated by using quarkcoin algorithm

algorithm.h

@@ -28,7 +28,8 @@ typedef enum {
   ALGO_NEOSCRYPT,
   ALGO_LYRA2RE,
   ALGO_PLUCK,
-  ALGO_YESCRYPT
+  ALGO_YESCRYPT,
+  ALGO_YESCRYPT_MULTI,
 } algorithm_type_t;
 
 extern const char *algorithm_type_str[];

driver-opencl.c

@@ -1431,7 +1431,7 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
       status = clEnqueueNDRangeKernel(clState->commandQueue, clState->extra_kernels[i], 1, p_global_work_offset,
                       globalThreads, localThreads, 0,  NULL, NULL);
       if (unlikely(status != CL_SUCCESS)) {
-          applog(LOG_ERR, "Error %d: Enqueueing kernel onto command queue. (clEnqueueNDRangeKernel)", status);
+          applog(LOG_ERR, "Error %d: Enqueueing kernel onto command queue. (clEnqueueNDRangeKernel)  %d", status,i);
           return -1;
       }
   }
@@ -1483,6 +1483,12 @@ static void opencl_thread_shutdown(struct thr_info *thr)
     clFinish(clState->commandQueue);
     clReleaseMemObject(clState->outputBuffer);
     clReleaseMemObject(clState->CLbuffer0);
+	if (clState->buffer1)
+	clReleaseMemObject(clState->buffer1);
+	if (clState->buffer2)
+	clReleaseMemObject(clState->buffer2);
+	if (clState->buffer3)
+	clReleaseMemObject(clState->buffer3);
     if (clState->padbuffer8)
       clReleaseMemObject(clState->padbuffer8);
     clReleaseKernel(clState->kernel);

example.bat

@@ -6,7 +6,7 @@ del *.bin
 @rem sgminer.exe --no-submit-stale --kernel Lyra2RE -o stratum+tcp://pool.verters.com:4444 -u djm34t.user -p password --gpu-platform 2 
 @rem sgminer.exe --no-submit-stale --kernel pluck -o stratum+tcp://sup.suprnova.cc:7777  -u djm34.2 -p password   --gpu-platform 2 --thread-concurrency 8192 -w 4  -I 12
 @rem sgminer.exe --no-submit-stale --kernel yescrypt -o stratum+tcp://mine2.bsty.nonce-pool.com:4095 -u djm34.1 -p password  --gpu-platform 1 -w 32 --thread-concurrency 512 --text-only --debug
-@rem sgminer.exe --no-submit-stale --kernel yescrypt -o stratum+tcp://mine2.bsty.nonce-pool.com:4095 -u djm34.1 -p password  --gpu-platform 1 -w 32 --thread-concurrency 512 --text-only -D
+sgminer.exe --no-submit-stale --kernel yescrypt -o stratum+tcp://mine2.bsty.nonce-pool.com:4095 -u djm34.1 -p password  --gpu-platform 1 -w 8 --thread-concurrency 1024 -I 9 
 
-sgminer.exe --no-submit-stale --kernel yescrypt -o stratum+tcp://mine2.bsty.nonce-pool.com:4095 -u djm34.1 -p password  --gpu-platform 0  -w 16 -g 2
+@rem sgminer.exe --no-submit-stale --kernel yescrypt-multi -o stratum+tcp://mine2.bsty.nonce-pool.com:4095 -u djm34.1 -p password  --gpu-platform 0 -w 4   -g 2 
 pause

kernel/yescrypt-multi.cl

@@ -0,0 +1,314 @@
+/*
+* "yescrypt" kernel implementation.
+*
+* ==========================(LICENSE BEGIN)============================
+*
+* Copyright (c) 2015  djm34
+*
+* Permission is hereby granted, free of charge, to any person obtaining
+* a copy of this software and associated documentation files (the
+* "Software"), to deal in the Software without restriction, including
+* without limitation the rights to use, copy, modify, merge, publish,
+* distribute, sublicense, and/or sell copies of the Software, and to
+* permit persons to whom the Software is furnished to do so, subject to
+* the following conditions:
+*
+* The above copyright notice and this permission notice shall be
+* included in all copies or substantial portions of the Software.
+*
+* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+* EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+* MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+* IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+* CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+* TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+* SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
+*
+* ===========================(LICENSE END)=============================
+*
+* @author   djm34
+*/
+#if !defined(cl_khr_byte_addressable_store)
+#error "Device does not support unaligned stores"
+#endif
+
+#include "yescrypt_essential.cl"
+
+__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
+__kernel void search(__global const uchar* restrict input, __global uint* restrict output, __global uchar *padcache, __global uchar* buff1, __global uchar* buff2, __global uchar* buff3, const uint target)
+{
+
+	__global ulong16 *hashbuffer = (__global ulong16 *)(padcache + (2048 * 128 * sizeof(ulong)* (get_global_id(0) % MAX_GLOBAL_THREADS)));
+	__global ulong16 *prevstate = (__global ulong16 *)(buff1 + (64 * 128 * sizeof(ulong)*(get_global_id(0) % MAX_GLOBAL_THREADS)));
+	__global uint8 *sha256tokeep = (__global uint8 *)(buff3 + (8 * sizeof(uint)*(get_global_id(0) % MAX_GLOBAL_THREADS)));
+	__global ulong16 *Bdev = (__global ulong16 *)(buff2 + (8 * 128 * sizeof(ulong)* (get_global_id(0) % MAX_GLOBAL_THREADS)));
+
+
+
+	uint nonce = (get_global_id(0));
+	uint  data[20];
+	uint16 in;
+	uint8 state1, state2;
+//	uint8 sha256tokeep;
+
+//	ulong16 Bdev[8]; // will require an additional buffer
+	((uint16 *)data)[0] = ((__global const uint16 *)input)[0];
+	((uint4  *)data)[4] = ((__global const uint4  *)input)[4];
+//	for (int i = 0; i<20; i++) { data[i] = SWAP32(data[i]); }
+	//	if (nonce == 10) { printf("data %08x %08x\n", data[0], data[1]); }
+	uint8 passwd = sha256_80(data, nonce);
+	//pbkdf		
+	in.lo = pad1.lo ^ passwd;
+	in.hi = pad1.hi;
+	state1 = sha256_Transform(in, H256);
+
+	in.lo = pad2.lo ^ passwd;
+	in.hi = pad2.hi;
+	state2 = sha256_Transform(in, H256);
+
+	in = ((uint16*)data)[0];
+	state1 = sha256_Transform(in, state1);
+#pragma unroll 1
+	for (int i = 0; i<8; i++)
+	{
+		uint16 result;
+		in = pad3;
+		in.s0 = data[16];
+		in.s1 = data[17];
+		in.s2 = data[18];
+		in.s3 = nonce;
+		in.s4 = 4 * i + 1;
+		in.lo = sha256_Transform(in, state1);
+		in.hi = pad4;
+		result.lo = swapvec(sha256_Transform(in, state2));
+		if (i == 0) sha256tokeep[0] = result.lo;
+		in = pad3;
+		in.s0 = data[16];
+		in.s1 = data[17];
+		in.s2 = data[18];
+		in.s3 = nonce;
+		in.s4 = 4 * i + 2;
+		in.lo = sha256_Transform(in, state1);
+		in.hi = pad4;
+		result.hi = swapvec(sha256_Transform(in, state2));
+		Bdev[i].lo = as_ulong8(shuffle(result));
+//		Bdev[i].lo = as_ulong8(result);
+		in = pad3;
+		in.s0 = data[16];
+		in.s1 = data[17];
+		in.s2 = data[18];
+		in.s3 = nonce;
+		in.s4 = 4 * i + 3;
+		in.lo = sha256_Transform(in, state1);
+		in.hi = pad4;
+		result.lo = swapvec(sha256_Transform(in, state2));
+		in = pad3;
+		in.s0 = data[16];
+		in.s1 = data[17];
+		in.s2 = data[18];
+		in.s3 = nonce;
+		in.s4 = 4 * i + 4;
+		in.lo = sha256_Transform(in, state1);
+		in.hi = pad4;
+		result.hi = swapvec(sha256_Transform(in, state2));
+
+
+		Bdev[i].hi = as_ulong8(shuffle(result));
+//		Bdev[i].hi = as_ulong8(result);
+	}
+
+	//mixing1
+
+	prevstate[0] = Bdev[0];
+	Bdev[0] = blockmix_salsa8_small2(Bdev[0]);
+	prevstate[1] = Bdev[0];
+	Bdev[0] = blockmix_salsa8_small2(Bdev[0]);
+
+	uint n = 1;
+#pragma unroll 1
+	for (uint i = 2; i < 64; i++)
+	{
+
+		prevstate[i] = Bdev[0];
+
+		if ((i&(i - 1)) == 0) n = n << 1;
+
+		uint j = as_uint2(Bdev[0].hi.s0).x & (n - 1);
+
+		j += i - n;
+		Bdev[0] ^= prevstate[j];
+
+		Bdev[0] = blockmix_salsa8_small2(Bdev[0]);
+	}
+
+
+}
+
+
+
+__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
+__kernel void search1(__global uchar *buffer1, __global uchar *buffer2)
+{
+}
+
+
+
+
+__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
+__kernel void search2(__global uchar *padcache, __global uchar *buff1, __global uchar *buff2)
+{
+
+			__global ulong16 *hashbuffer = (__global ulong16 *)(padcache + (2048 * 128 * sizeof(ulong)* (get_global_id(0) % MAX_GLOBAL_THREADS)));
+			__global ulong16* prevstate =  (__global ulong16 *)(buff1 +      (64 * 128 * sizeof(ulong)* (get_global_id(0) % MAX_GLOBAL_THREADS)));
+			__global ulong16 *Bdev =       (__global ulong16 *)(buff2 +       (8 * 128 * sizeof(ulong)* (get_global_id(0) % MAX_GLOBAL_THREADS)));
+
+
+			for (int i = 0; i<8; i++)
+				hashbuffer[i] = Bdev[i];
+
+			blockmix_pwxform((__global ulong8*)Bdev, prevstate);
+
+
+			for (int i = 0; i<8; i++)
+				hashbuffer[i + 8] = Bdev[i];
+
+			blockmix_pwxform((__global ulong8*)Bdev, prevstate);
+			int n = 1;
+#pragma unroll 1
+			for (int i = 2; i < 2048; i ++)
+			{
+
+				for (int k = 0; k<8; k++)
+					(hashbuffer + 8 * i)[k] = Bdev[k];
+
+
+				if ((i&(i - 1)) == 0) n = n << 1;
+
+				uint j = as_uint2(Bdev[7].hi.s0).x & (n - 1);
+				j += i - n;
+
+				for (int k = 0; k < 8; k++)
+					Bdev[k] ^= (hashbuffer + 8 * j)[k];
+
+
+				blockmix_pwxform((__global ulong8*)Bdev, prevstate);
+            }
+}
+
+/*
+__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
+__kernel void search3(__global uchar *buffer1, __global uchar *buffer2)
+{
+}
+*/
+
+
+__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
+__kernel void search3(__global uchar *padcache, __global uchar *buff1, __global uchar *buff2)
+{
+
+	__global ulong16 *hashbuffer = (__global ulong16 *)(padcache + (2048 * 128 * sizeof(ulong)* (get_global_id(0) % MAX_GLOBAL_THREADS)));
+	__global ulong16* prevstate = (__global ulong16 *)(buff1 + (64 * 128 * sizeof(ulong)* (get_global_id(0) % MAX_GLOBAL_THREADS)));
+	__global ulong16 *Bdev = (__global ulong16 *)(buff2 + (8 * 128 * sizeof(ulong)* (get_global_id(0) % MAX_GLOBAL_THREADS)));
+
+
+#pragma unroll 1
+	for (int z = 0; z < 684; z++)
+	{
+
+		uint j = as_uint2(Bdev[7].hi.s0).x & 2047;
+
+
+		for (int k = 0; k < 8; k++)
+			Bdev[k] ^= (hashbuffer + 8 * j)[k];
+
+		if (z<682)
+			for (int k = 0; k<8; k++)
+				(hashbuffer + 8 * j)[k] = Bdev[k];
+
+		blockmix_pwxform((__global ulong8*)Bdev, prevstate);
+////
+	}
+
+}
+
+/*
+__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
+__kernel void search5(__global uchar *buffer1, __global uchar *buffer2)
+{
+}
+*/
+
+__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
+__kernel void search4(__global const uchar* restrict input, __global uint* restrict output, __global uchar *buff2,__global uchar* buff3, const uint target)
+{
+
+	__global ulong16 *Bdev =     (__global ulong16 *)(buff2 + (8 * 128 * sizeof(ulong)* (get_global_id(0) % MAX_GLOBAL_THREADS)));
+	__global uint8 *sha256tokeep = (__global uint8 *)(buff3 + (8 * sizeof(uint)*(get_global_id(0) % MAX_GLOBAL_THREADS)));
+
+	uint nonce = (get_global_id(0));
+
+
+	uint  data[20];
+	((uint16 *)data)[0] = ((__global const uint16 *)input)[0];
+	((uint4  *)data)[4] = ((__global const uint4  *)input)[4];
+//	for (int i = 0; i<20; i++) { data[i] = SWAP32(data[i]); }
+	uint8 swpass = swapvec(sha256tokeep[0]);
+    uint16 in;
+    uint8 state1,state2;
+	in.lo = pad1.lo ^ swpass;
+	in.hi = pad1.hi;
+
+
+	state1 = sha256_Transform(in, H256);
+
+	in.lo = pad2.lo ^ swpass;
+	in.hi = pad2.hi;
+	state2 = sha256_Transform(in, H256);
+
+#pragma unroll 1
+	for (int i = 0; i<8; i++) {
+		in = unshuffle(Bdev[i].lo);
+		in = swapvec16(in);
+		state1 = sha256_Transform(in, state1);
+		in = unshuffle(Bdev[i].hi);
+		in = swapvec16(in);
+		state1 = sha256_Transform(in, state1);
+	}
+	in = pad5;
+	state1 = sha256_Transform(in, state1);
+	in.lo = state1;
+	in.hi = pad4;
+	uint8 res = sha256_Transform(in, state2);
+
+	//hmac and final sha
+
+	in.lo = pad1.lo ^ res;
+	in.hi = pad1.hi;
+	state1 = sha256_Transform(in, H256);
+	in.lo = pad2.lo ^ res;
+	in.hi = pad2.hi;
+	state2 = sha256_Transform(in, H256);
+	in = ((uint16*)data)[0];
+	state1 = sha256_Transform(in, state1);
+	in = padsha80;
+	in.s0 = data[16];
+	in.s1 = data[17];
+	in.s2 = data[18];
+	in.s3 = get_global_id(0);
+	in.sf = 0x480;
+	state1 = sha256_Transform(in, state1);
+	in.lo = state1;
+	in.hi = pad4;
+	state1 = sha256_Transform(in, state2);
+	//		state2 = H256;
+	in.lo = state1;
+	in.hi = pad4;
+	in.sf = 0x100;
+	res = sha256_Transform(in, H256);
+
+
+	if (SWAP32(res.s7) <= (target)) 
+		output[atomic_inc(output + 0xFF)] = (nonce);
+	
+}

kernel/yescrypt.cl

@@ -35,237 +35,219 @@
 #include "yescrypt_essential.cl"
 
 
-
-
 __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
-__kernel void search(__global const uchar* restrict input, __global uint* restrict output, __global uchar *padcache, __global uchar* buff2, const uint target)
+__kernel void search(__global const uchar* restrict input, __global uint* restrict output, __global uchar *padcache, __global uchar* buff1, __global uchar* buff2, const uint target)
 {
 
-	__global ulong16 *hashbuffer = (__global ulong16 *)(padcache + (2048*128 *sizeof(ulong)* (get_global_id(0) % MAX_GLOBAL_THREADS)));
+	__global ulong16 *hashbuffer = (__global ulong16 *)(padcache + (2048 * 128 * sizeof(ulong)* (get_global_id(0) % MAX_GLOBAL_THREADS)));
-	__global ulong16 *prevstate = (__global ulong16 *)(buff2 + (64 * 128 * sizeof(ulong)*(get_global_id(0) % MAX_GLOBAL_THREADS)));
+	__global ulong16 *prevstate = (__global ulong16 *)(buff1 + (64 * 128 * sizeof(ulong)*(get_global_id(0) % MAX_GLOBAL_THREADS)));
+	__global ulong16 *Bdev = (__global ulong16 *)(buff2 + (8 * 128 * sizeof(ulong)* (get_global_id(0) % MAX_GLOBAL_THREADS)));
 
 
 
-    uint nonce = (get_global_id(0));
+	uint nonce = (get_global_id(0));
-    uint  data[20];
+	uint  data[20];
-    uint16 in;
+	uint16 in;
-    uint8 state1, state2;
+	uint8 state1, state2;
     uint8 sha256tokeep;
- 
+
-    ulong16 Bdev[8]; // will require an additional buffer
 	((uint16 *)data)[0] = ((__global const uint16 *)input)[0];
 	((uint4  *)data)[4] = ((__global const uint4  *)input)[4];
 	for (int i = 0; i<20; i++) { data[i] = SWAP32(data[i]); }
-//	if (nonce == 10) { printf("data %08x %08x\n", data[0], data[1]); }
+	//	if (nonce == 10) { printf("data %08x %08x\n", data[0], data[1]); }
 	uint8 passwd = sha256_80(data, nonce);
-//pbkdf		
+	//pbkdf		
-		in.lo = pad1.lo ^ passwd;
+	in.lo = pad1.lo ^ passwd;
-		in.hi = pad1.hi;
+	in.hi = pad1.hi;
-		state1 = sha256_Transform(in, H256);
+	state1 = sha256_Transform(in, H256);
 
-		in.lo = pad2.lo ^ passwd;
+	in.lo = pad2.lo ^ passwd;
-		in.hi = pad2.hi;
+	in.hi = pad2.hi;
-		state2 = sha256_Transform(in, H256);
+	state2 = sha256_Transform(in, H256);
 
-		in = ((uint16*)data)[0];
+	in = ((uint16*)data)[0];
-		state1 = sha256_Transform(in, state1);
+	state1 = sha256_Transform(in, state1);
- 
-		for (int i = 0; i<8; i++)
-		{
-			uint16 result;
-			in = pad3;
-			in.s0 = data[16];
-			in.s1 = data[17];
-			in.s2 = data[18];
-			in.s3 = nonce;
-            in.s4 = 4*i+1;
-			in.lo = sha256_Transform(in, state1);
-			in.hi = pad4;
-			result.lo = swapvec(sha256_Transform(in, state2));
-			if (i == 0) sha256tokeep = result.lo;
-			in = pad3;
-			in.s0 = data[16];
-			in.s1 = data[17];
-			in.s2 = data[18];
-			in.s3 = nonce;
-			in.s4 = 4 * i + 2;
-			in.lo = sha256_Transform(in, state1); 			
-			in.hi = pad4;
-			result.hi = swapvec(sha256_Transform(in, state2));
-            Bdev[i].lo = as_ulong8(shuffle(result));
-			in = pad3;
-			in.s0 = data[16];
-			in.s1 = data[17];
-			in.s2 = data[18];
-			in.s3 = nonce;
-			in.s4 = 4 * i + 3;
-			in.lo = sha256_Transform(in, state1);
-			in.hi = pad4;
-			result.lo = swapvec(sha256_Transform(in, state2));
-			in = pad3;
-			in.s0 = data[16];
-			in.s1 = data[17];
-			in.s2 = data[18];
-			in.s3 = nonce;
-			in.s4 =  4 * i + 4;
-			in.lo = sha256_Transform(in, state1);
-			in.hi = pad4;
-			result.hi = swapvec(sha256_Transform(in, state2));
-
-			
-			Bdev[i].hi = as_ulong8(shuffle(result));
-		}
-
-//mixing1
-		
-		prevstate[0] = Bdev[0];
-		Bdev[0]=blockmix_salsa8_small2(Bdev[0]);
-		prevstate[1] = Bdev[0];
-		Bdev[0] = blockmix_salsa8_small2(Bdev[0]);
-
-
-
-		uint n = 1;
 #pragma unroll 1
-		for (uint i = 2; i < 64; i++)
+	for (int i = 0; i<8; i++)
-		{
+	{
-
+		uint16 result;
-			prevstate[i] = Bdev[0];
+		in = pad3;
-
-			if ((i&(i - 1)) == 0) n = n << 1;
-
-			uint j = as_uint2(Bdev[0].hi.s0).x & (n - 1);
-
-			j += i - n;
-			Bdev[0] ^= prevstate[j];
-
-			Bdev[0] = blockmix_salsa8_small2(Bdev[0]);
-		}
-//mixing1_2
-
-		for (int i = 0; i<8; i++)
-			hashbuffer[i] = Bdev[i];
-
-		blockmix_pwxform((ulong8*)Bdev,prevstate);
-
-
-		for (int i = 0; i<8; i++)
-			hashbuffer[i + 8] = Bdev[i];
-
-		blockmix_pwxform((ulong8*)Bdev,prevstate);
-		n = 1;
-#pragma unroll 1
-		for (int i = 2; i < 2048; i++)
-		{
-
-
-			for (int k = 0; k<8; k++)
-				(hashbuffer + 8 * i)[k] = Bdev[k];
-
-
-			if ((i&(i - 1)) == 0) n = n << 1;
-
-			uint j = as_uint2(Bdev[7].hi.s0).x & (n - 1);
-			j += i - n;
-
-			for (int k = 0; k < 8; k++)
-				Bdev[k] ^= (hashbuffer + 8 * j)[k];
-
-
-			blockmix_pwxform( (ulong8*)Bdev,prevstate);
-
-		}
-/////////////////////////
-
-////mix2_2
-
-
-		//#pragma unroll
-#pragma unroll 1
-		for (int z = 0; z < 684; z++)
-		{
-
-			uint j = as_uint2(Bdev[7].hi.s0).x & 2047;
-
-
-			for (int k = 0; k < 8; k++)
-				Bdev[k] ^= (hashbuffer + 8 * j)[k];
-
-			if (z<682)
-				for (int k = 0; k<8; k++)
-					(hashbuffer+8 * j)[k] = Bdev[k];
-
-			blockmix_pwxform((ulong8*)Bdev,prevstate);
-
-		}
-
-
-		for (int i = 0; i<8; i++) {
-			Bdev[i].lo = as_ulong8(unshuffle(Bdev[i].lo));
-			Bdev[i].hi = as_ulong8(unshuffle(Bdev[i].hi));
-		}
-/////////////////////////////////////
-///////// pbkdf final
-
-
-		uint8 swpass = swapvec(sha256tokeep);
-		in.lo = pad1.lo ^ swpass;
-		in.hi = pad1.hi;
-
-		state1 = sha256_Transform(in, H256);
-
-		in.lo = pad2.lo ^ swpass;
-		in.hi = pad2.hi;
-		state2 = sha256_Transform(in, H256);
-
-		for (int i = 0; i<8; i++) {
-			in = as_uint16(Bdev[i].lo);
-			in = swapvec16(in);
-			state1 = sha256_Transform(in, state1);
-			in = as_uint16(Bdev[i].hi);
-			in = swapvec16(in);
-			state1 = sha256_Transform(in, state1);
-		}
-		in = pad5;
-		state1 = sha256_Transform(in, state1);
-		in.lo = state1;
-		in.hi = pad4;
-		uint8 res = sha256_Transform(in, state2);
-
-		//hmac and final sha
-
-		in.lo = pad1.lo ^ res;
-		in.hi = pad1.hi;
-		state1 = sha256_Transform(in, H256);
-		in.lo = pad2.lo ^ res;
-		in.hi = pad2.hi;
-		state2 = sha256_Transform(in, H256);
-		in = ((uint16*)data)[0];
-		state1 = sha256_Transform(in, state1);
-		in = padsha80;
 		in.s0 = data[16];
 		in.s1 = data[17];
 		in.s2 = data[18];
 		in.s3 = nonce;
-		in.sf = 0x480;
+		in.s4 = 4 * i + 1;
+		in.lo = sha256_Transform(in, state1);
+		in.hi = pad4;
+		result.lo = swapvec(sha256_Transform(in, state2));
+		if (i == 0) sha256tokeep = result.lo;
+		in = pad3;
+		in.s0 = data[16];
+		in.s1 = data[17];
+		in.s2 = data[18];
+		in.s3 = nonce;
+		in.s4 = 4 * i + 2;
+		in.lo = sha256_Transform(in, state1);
+		in.hi = pad4;
+		result.hi = swapvec(sha256_Transform(in, state2));
+		Bdev[i].lo = as_ulong8(shuffle(result));
+		in = pad3;
+		in.s0 = data[16];
+		in.s1 = data[17];
+		in.s2 = data[18];
+		in.s3 = nonce;
+		in.s4 = 4 * i + 3;
+		in.lo = sha256_Transform(in, state1);
+		in.hi = pad4;
+		result.lo = swapvec(sha256_Transform(in, state2));
+		in = pad3;
+		in.s0 = data[16];
+		in.s1 = data[17];
+		in.s2 = data[18];
+		in.s3 = nonce;
+		in.s4 = 4 * i + 4;
+		in.lo = sha256_Transform(in, state1);
+		in.hi = pad4;
+		result.hi = swapvec(sha256_Transform(in, state2));
+
+
+		Bdev[i].hi = as_ulong8(shuffle(result));
+	}
+
+	//mixing1
+
+	prevstate[0] = Bdev[0];
+	Bdev[0] = blockmix_salsa8_small2(Bdev[0]);
+	prevstate[1] = Bdev[0];
+	Bdev[0] = blockmix_salsa8_small2(Bdev[0]);
+
+	uint n = 1;
+#pragma unroll 1
+	for (uint i = 2; i < 64; i++)
+	{
+
+		prevstate[i] = Bdev[0];
+
+		if ((i&(i - 1)) == 0) n = n << 1;
+
+		uint j = as_uint2(Bdev[0].hi.s0).x & (n - 1);
+
+		j += i - n;
+		Bdev[0] ^= prevstate[j];
+
+		Bdev[0] = blockmix_salsa8_small2(Bdev[0]);
+	}
+
+
+	for (int i = 0; i<8; i++)
+		hashbuffer[i] = Bdev[i];
+
+	blockmix_pwxform((__global ulong8*)Bdev, prevstate);
+
+
+	for (int i = 0; i<8; i++)
+		hashbuffer[i + 8] = Bdev[i];
+
+	blockmix_pwxform((__global ulong8*)Bdev, prevstate);
+    n = 1;
+#pragma unroll 1
+	for (int i = 2; i < 2048; i++)
+	{
+
+		for (int k = 0; k<8; k++)
+			(hashbuffer + 8 * i)[k] = Bdev[k];
+
+
+		if ((i&(i - 1)) == 0) n = n << 1;
+
+		uint j = as_uint2(Bdev[7].hi.s0).x & (n - 1);
+		j += i - n;
+
+		for (int k = 0; k < 8; k++)
+			Bdev[k] ^= (hashbuffer + 8 * j)[k];
+
+
+		blockmix_pwxform((__global ulong8*)Bdev, prevstate);
+	}
+
+
+#pragma unroll 1
+	for (int z = 0; z < 684; z++)
+	{
+
+		uint j = as_uint2(Bdev[7].hi.s0).x & 2047;
+
+
+		for (int k = 0; k < 8; k++)
+			Bdev[k] ^= (hashbuffer + 8 * j)[k];
+
+		if (z<682)
+			for (int k = 0; k<8; k++)
+				(hashbuffer + 8 * j)[k] = Bdev[k];
+
+		blockmix_pwxform((__global ulong8*)Bdev, prevstate);
+		////
+	}
+
+
+
+	uint8 swpass = swapvec(sha256tokeep);
+//	uint16 in;
+//	uint8 state1, state2;
+	in.lo = pad1.lo ^ swpass;
+	in.hi = pad1.hi;
+
+
+	state1 = sha256_Transform(in, H256);
+
+	in.lo = pad2.lo ^ swpass;
+	in.hi = pad2.hi;
+	state2 = sha256_Transform(in, H256);
+
+#pragma unroll 1
+	for (int i = 0; i<8; i++) {
+		in = unshuffle(Bdev[i].lo);
+		in = swapvec16(in);
 		state1 = sha256_Transform(in, state1);
-		in.lo = state1;
+		in = unshuffle(Bdev[i].hi);
-		in.hi = pad4;
+		in = swapvec16(in);
-		state1 = sha256_Transform(in, state2);
+		state1 = sha256_Transform(in, state1);
-		//		state2 = H256;
+	}
-		in.lo = state1;
+	in = pad5;
-		in.hi = pad4;
+	state1 = sha256_Transform(in, state1);
-		in.sf = 0x100;
+	in.lo = state1;
-		res = sha256_Transform(in, H256);
+	in.hi = pad4;
-		//		return(swapvec(res));
+	uint8 res = sha256_Transform(in, state2);
+
+	//hmac and final sha
+
+	in.lo = pad1.lo ^ res;
+	in.hi = pad1.hi;
+	state1 = sha256_Transform(in, H256);
+	in.lo = pad2.lo ^ res;
+	in.hi = pad2.hi;
+	state2 = sha256_Transform(in, H256);
+	in = ((uint16*)data)[0];
+	state1 = sha256_Transform(in, state1);
+	in = padsha80;
+	in.s0 = data[16];
+	in.s1 = data[17];
+	in.s2 = data[18];
+	in.s3 = get_global_id(0);
+	in.sf = 0x480;
+	state1 = sha256_Transform(in, state1);
+	in.lo = state1;
+	in.hi = pad4;
+	state1 = sha256_Transform(in, state2);
+	//		state2 = H256;
+	in.lo = state1;
+	in.hi = pad4;
+	in.sf = 0x100;
+	res = sha256_Transform(in, H256);
 
 
-//		if (nonce == 10) {  }
+	if (SWAP32(res.s7) <= (target))
+		output[atomic_inc(output + 0xFF)] = (nonce);
 
-		if ( SWAP32(res.s7) <= (target)) {
-			output[atomic_inc(output + 0xFF)] = (nonce);
-			//printf("gpu hashbuffer %08x nonce %08x\n",((__global uint *)hashbuffer)[7] ,SWAP32(get_global_id(0)));    		
 }
 
-
-}

kernel/yescrypt_essential.cl

@@ -140,6 +140,8 @@ inline uint8 swapvec(uint8 buf)
 	return vec;
 }
 
+
+
 inline uint16 swapvec16(uint16 buf)
 {
 	uint16 vec;
@@ -173,6 +175,18 @@ uint t;
 	return(as_ulong8(st + Bx));
 }
 
+ ulong8 salsa20_8n(uint16 Bx)
+ {
+	 uint t;
+	 uint16 st = Bx;
+	 SALSA_CORE(st);
+	 SALSA_CORE(st);
+	 SALSA_CORE(st);
+	 SALSA_CORE(st);
+	 return(as_ulong8(st + Bx));
+ }
+
+
  ulong16 blockmix_salsa8_small2(ulong16 Bin)
 {
 	ulong8 X = Bin.hi;
@@ -346,7 +360,7 @@ inline ulong8 block_pwxform_long(ulong8 Bout, __global ulong2 *prevstate)
 
 
 
-inline void blockmix_pwxform( ulong8 *Bin, __global  ulong16 *prevstate)
+inline void blockmix_pwxform(__global ulong8 *Bin, __global  ulong16 *prevstate)
 {
 	Bin[0] ^= Bin[15];
 	Bin[0] = block_pwxform_long_old(Bin[0], prevstate);

ocl.c

@@ -516,7 +516,8 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize, algorithm_t *alg
 
 	  applog(LOG_DEBUG, "GPU %d: computing max. global thread count to %u", gpu, (unsigned)(cgpu->thread_concurrency));
    }
-   else if (!safe_cmp(cgpu->algorithm.name, "yescrypt") && !cgpu->opt_tc) {
+  else if ((!safe_cmp(cgpu->algorithm.name, "yescrypt") ||
+            !safe_cmp(algorithm->name, "yescrypt-multi")) && !cgpu->opt_tc) {
 		  size_t glob_thread_count;
 		  long max_int;
 		  unsigned char type = 0;
@@ -702,6 +703,12 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize, algorithm_t *alg
   }
 
   size_t bufsize;
+  size_t buf1size;
+  size_t buf3size;
+  size_t buf2size;
+
+
+
   size_t readbufsize = 128;
 
   if (algorithm->rw_buffer_size < 0) {
@@ -717,10 +724,12 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize, algorithm_t *alg
       applog(LOG_DEBUG, "Neoscrypt buffer sizes: %lu RW, %lu R", (unsigned long)bufsize, (unsigned long)readbufsize);
     // scrypt/n-scrypt
     } 
-	else if (!safe_cmp(algorithm->name, "yescrypt")) {
+	else if (!safe_cmp(algorithm->name, "yescrypt") || !safe_cmp(algorithm->name, "yescrypt-multi")) {
 		/* The scratch/pad-buffer needs 32kBytes memory per thread. */
 		bufsize = YESCRYPT_SCRATCHBUF_SIZE * cgpu->thread_concurrency;
-
+		buf1size = PLUCK_SECBUF_SIZE * cgpu->thread_concurrency;
+		buf2size = 128 * 8 * 8 * cgpu->thread_concurrency;
+		buf3size= 8 * 8 * 4 * cgpu->thread_concurrency;
 		/* This is the input buffer. For yescrypt this is guaranteed to be
 		* 80 bytes only. */
 		readbufsize = 80;
@@ -753,6 +762,10 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize, algorithm_t *alg
   }
 
   clState->padbuffer8 = NULL;
+  clState->buffer1 = NULL;
+  clState->buffer2 = NULL;
+  clState->buffer3 = NULL;
+
 
   if (bufsize > 0) {
     applog(LOG_DEBUG, "Creating read/write buffer sized %lu", (unsigned long)bufsize);
@@ -765,20 +778,20 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize, algorithm_t *alg
     }
 
 
-	if (!safe_cmp(algorithm->name, "yescrypt")) {
+	if (!safe_cmp(algorithm->name, "yescrypt") || !safe_cmp(algorithm->name, "yescrypt-multi")) {
 // need additionnal buffers
-		clState->buffer1 = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, PLUCK_SECBUF_SIZE * cgpu->thread_concurrency, NULL, &status);
+		clState->buffer1 = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, buf1size, NULL, &status);
 		if (status != CL_SUCCESS && !clState->buffer1) {
 			applog(LOG_DEBUG, "Error %d: clCreateBuffer (buffer1), decrease TC or increase LG", status);
 			return NULL;}
 
-		clState->buffer2 = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, 128 * 8 * 8 * cgpu->thread_concurrency, NULL, &status);
+		clState->buffer2 = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, buf2size, NULL, &status);
 		if (status != CL_SUCCESS && !clState->buffer2) {
 			applog(LOG_DEBUG, "Error %d: clCreateBuffer (buffer2), decrease TC or increase LG", status);
 			return NULL;
 		}
 
-		clState->buffer3 = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, 8 * 8 * 4 * cgpu->thread_concurrency, NULL, &status);
+		clState->buffer3 = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, buf3size, NULL, &status);
 		if (status != CL_SUCCESS && !clState->buffer3) {
 			applog(LOG_DEBUG, "Error %d: clCreateBuffer (buffer3), decrease TC or increase LG", status);
 			return NULL;

sgminer.c

@@ -7082,7 +7082,9 @@ bool test_nonce(struct work *work, uint32_t nonce)
 
   // for Neoscrypt, the diff1targ value is in work->target
   if (!safe_cmp(work->pool->algorithm.name, "neoscrypt") || !safe_cmp(work->pool->algorithm.name, "pluck")
-	  || !safe_cmp(work->pool->algorithm.name, "yescrypt") ) {
+	  || !safe_cmp(work->pool->algorithm.name, "yescrypt") 
+	  || !safe_cmp(work->pool->algorithm.name, "yescrypt-multi")
+) {
     diff1targ = ((uint32_t *)work->target)[7];
   }
   else {