Merge commit '537b28d' into scrypt

driver-cpu.h

@@ -60,5 +60,6 @@ extern void show_algo(char buf[OPT_SHOW_LEN], const enum sha256_algos *algo);
 extern char *force_nthreads_int(const char *arg, int *i);
 extern void init_max_name_len();
 extern double bench_algo_stage3(enum sha256_algos algo);
+extern void set_scrypt_algo(enum sha256_algos *algo);
 
 #endif /* __DEVICE_CPU_H__ */

driver-opencl.c

@@ -1210,27 +1210,22 @@ static cl_int queue_diablo_kernel(_clState *clState, dev_blk_ctx *blk, cl_uint t
 }
 
 #ifdef USE_SCRYPT
-static cl_int queue_scrypt_kernel(_clState *clState, dev_blk_ctx *blk, cl_uint threads)
+static cl_int queue_scrypt_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_unused cl_uint threads)
 {
-	cl_uint4 *midstate = (cl_uint4 *)blk->midstate;
+	char *midstate = blk->work->midstate;
 	cl_kernel *kernel = &clState->kernel;
 	unsigned int num = 0;
 	cl_int status = 0;
-	int i;
+
+	clState->cldata = blk->work->data;
+	status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 80, clState->cldata, 0, NULL,NULL);
 
 	CL_SET_ARG(clState->CLbuffer0);
 	CL_SET_ARG(clState->outputBuffer);
 	CL_SET_ARG(clState->padbuffer8);
-	CL_SET_ARG(midstate[0]);
-	CL_SET_ARG(midstate[16]);
-
-#if 0
-	clSetKernelArg(clState->kernel,0,sizeof(cl_mem), &clState->CLbuffer[0]);
-	clSetKernelArg(clState->kernel,1,sizeof(cl_mem), &clState->CLbuffer[1]);
-	clSetKernelArg(clState->kernel,2,sizeof(cl_mem), &clState->padbuffer8);
-	clSetKernelArg(clState->kernel,3,sizeof(cl_uint4), &midstate[0]);
-	clSetKernelArg(clState->kernel,4,sizeof(cl_uint4), &midstate[16]);
-#endif
+	CL_SET_VARG(4, &midstate[0]);
+	CL_SET_VARG(4, &midstate[16]);
+
 	return status;
 }
 #endif
@@ -1558,7 +1553,7 @@ static bool opencl_thread_init(struct thr_info *thr)
 	struct cgpu_info *gpu = thr->cgpu;
 	struct opencl_thread_data *thrdata;
 	_clState *clState = clStates[thr_id];
-	cl_int status;
+	cl_int status = 0;
 	thrdata = calloc(1, sizeof(*thrdata));
 	thr->cgpu_data = thrdata;
 
@@ -1596,11 +1591,7 @@ static bool opencl_thread_init(struct thr_info *thr)
 		return false;
 	}
 
-#ifdef USE_SCRYPT
-	if (opt_scrypt)
-		status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, BUFFERSIZE, blank_res, 0, NULL,NULL);
-#endif
-	status = clEnqueueWriteBuffer(clState->commandQueue, clState->outputBuffer, CL_TRUE, 0,
+	status |= clEnqueueWriteBuffer(clState->commandQueue, clState->outputBuffer, CL_TRUE, 0,
 			BUFFERSIZE, blank_res, 0, NULL, NULL);
 	if (unlikely(status != CL_SUCCESS)) {
 		applog(LOG_ERR, "Error: clEnqueueWriteBuffer failed.");
@@ -1629,7 +1620,12 @@ static void opencl_free_work(struct thr_info *thr, struct work *work)
 
 static bool opencl_prepare_work(struct thr_info __maybe_unused *thr, struct work *work)
 {
-	precalc_hash(&work->blk, (uint32_t *)(work->midstate), (uint32_t *)(work->data + 64));
+#ifdef USE_SCRYPT
+	if (opt_scrypt)
+		work->blk.work = work;
+	else
+#endif
+		precalc_hash(&work->blk, (uint32_t *)(work->midstate), (uint32_t *)(work->data + 64));
 	return true;
 }
 
@@ -1689,6 +1685,7 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
 			   localThreads[0], gpu->intensity);
 	if (hashes > gpu->max_hashes)
 		gpu->max_hashes = hashes;
+
 	status = thrdata->queue_kernel_parameters(clState, &work->blk, globalThreads[0]);
 	if (unlikely(status != CL_SUCCESS)) {
 		applog(LOG_ERR, "Error: clSetKernelArg of all params failed.");

findnonce.c

@@ -45,7 +45,8 @@ const uint32_t SHA256_K[64] = {
 	d = d + h; \
 	h = h + (rotate(a, 30) ^ rotate(a, 19) ^ rotate(a, 10)) + ((a & b) | (c & (a | b)))
 
-void precalc_hash(dev_blk_ctx *blk, uint32_t *state, uint32_t *data) {
+void precalc_hash(dev_blk_ctx *blk, uint32_t *state, uint32_t *data)
+{
 	cl_uint A, B, C, D, E, F, G, H;
 
 	A = state[0];
@@ -127,10 +128,6 @@ void precalc_hash(dev_blk_ctx *blk, uint32_t *state, uint32_t *data) {
 	blk->fiveA = blk->ctx_f + SHA256_K[5];
 	blk->sixA = blk->ctx_g + SHA256_K[6];
 	blk->sevenA = blk->ctx_h + SHA256_K[7];
-
-#ifdef USE_SCRYPT
-	blk->midstate = (unsigned char *)state;
-#endif
 }
 
 #define P(t) (W[(t)&0xF] = W[(t-16)&0xF] + (rotate(W[(t-15)&0xF], 25) ^ rotate(W[(t-15)&0xF], 14) ^ (W[(t-15)&0xF] >> 3)) + W[(t-7)&0xF] + (rotate(W[(t-2)&0xF], 15) ^ rotate(W[(t-2)&0xF], 13) ^ (W[(t-2)&0xF] >> 10)))
@@ -232,13 +229,16 @@ static void *postcalc_hash(void *userdata)
 	pthread_detach(pthread_self());
 
 	for (entry = 0; entry < FOUND; entry++) {
-		if (pcd->res[entry]) {
+		uint32_t nonce = pcd->res[entry];
+
+		if (nonce) {
+			applog(LOG_DEBUG, "OCL NONCE %u", nonce);
 #ifdef USE_SCRYPT
 			if (opt_scrypt)
-				submit_nonce(thr, pcd->work, pcd->res[entry]);
+				submit_nonce(thr, pcd->work, nonce);
 			else
 #endif
-				send_nonce(pcd, pcd->res[entry]);
+				send_nonce(pcd, nonce);
 		nonces++;
 		}
 	}

miner.c

@@ -1907,8 +1907,13 @@ static bool submit_upstream_work(const struct work *work, CURL *curl)
 
 	if (!QUIET) {
 		hash32 = (uint32_t *)(work->hash);
-		sprintf(hashshow, "%08lx.%08lx%s", (unsigned long)(hash32[6]), (unsigned long)(hash32[5]),
-			work->block? " BLOCK!" : "");
+		if (opt_scrypt) {
+			sprintf(hashshow, "%08lx.%08lx%s", (unsigned long)(hash32[7]), (unsigned long)(hash32[6]),
+				work->block? " BLOCK!" : "");
+		} else {
+			sprintf(hashshow, "%08lx.%08lx%s", (unsigned long)(hash32[6]), (unsigned long)(hash32[5]),
+				work->block? " BLOCK!" : "");
+		}
 	}
 
 	/* Theoretically threads could race when modifying accepted and
@@ -3157,6 +3162,11 @@ void write_config(FILE *fcfg)
 				case KL_DIABLO:
 					fprintf(fcfg, "diablo");
 					break;
+#ifdef USE_SCRYPT
+				case KL_SCRYPT:
+					fprintf(fcfg, "scrypt");
+					break;
+#endif
 			}
 		}
 #ifdef HAVE_ADL
@@ -4334,13 +4344,15 @@ bool hashtest(const struct work *work, bool checktarget)
 
 bool test_nonce(struct work *work, uint32_t nonce, bool checktarget)
 {
-	work->data[64 + 12 + 0] = (nonce >> 0) & 0xff;
-	work->data[64 + 12 + 1] = (nonce >> 8) & 0xff;
-	work->data[64 + 12 + 2] = (nonce >> 16) & 0xff;
-	work->data[64 + 12 + 3] = (nonce >> 24) & 0xff;
+	uint32_t *work_nonce = (uint32_t *)(work->data + 64 + 12);
 
-	if (opt_scrypt)
+	if (opt_scrypt) {
+		*work_nonce = nonce;
 		return true;
+	}
+
+	*work_nonce = htobe32(nonce);
+
 
 	return hashtest(work, checktarget);
 }

miner.h

@@ -679,7 +679,7 @@ typedef struct {
 	cl_uint zeroA, zeroB;
 	cl_uint oneA, twoA, threeA, fourA, fiveA, sixA, sevenA;
 #ifdef USE_SCRYPT
-	unsigned char *midstate;
+	struct work *work;
 #endif
 } dev_blk_ctx;
 #else

ocl.c

@@ -292,7 +292,7 @@ int clDevicesNum(void) {
 		status = clGetPlatformInfo(platform, CL_PLATFORM_VERSION, sizeof(pbuff), pbuff, NULL);
 		if (status == CL_SUCCESS)
 			applog(LOG_INFO, "CL Platform %d version: %s", i, pbuff);
-		status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &numDevices);
+		status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, NULL, &numDevices);
 		if (status != CL_SUCCESS) {
 			applog(LOG_ERR, "Error %d: Getting Device IDs (num)", status);
 			if (i < numPlatforms - 1)
@@ -309,7 +309,7 @@ int clDevicesNum(void) {
 			char pbuff[256];
 			cl_device_id *devices = (cl_device_id *)malloc(numDevices*sizeof(cl_device_id));
 
-			clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, numDevices, devices, NULL);
+			clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, numDevices, devices, NULL);
 			for (j = 0; j < numDevices; j++) {
 				clGetDeviceInfo(devices[j], CL_DEVICE_NAME, sizeof(pbuff), pbuff, NULL);
 				applog(LOG_INFO, "\t%i\t%s", j, pbuff);
@@ -432,7 +432,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
 	if (status == CL_SUCCESS)
 		applog(LOG_INFO, "CL Platform version: %s", vbuff);
 
-	status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &numDevices);
+	status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, NULL, &numDevices);
 	if (status != CL_SUCCESS) {
 		applog(LOG_ERR, "Error %d: Getting Device IDs (num)", status);
 		return NULL;
@@ -443,7 +443,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
 
 		/* Now, get the device list data */
 
-		status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, numDevices, devices, NULL);
+		status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, numDevices, devices, NULL);
 		if (status != CL_SUCCESS) {
 			applog(LOG_ERR, "Error %d: Getting Device IDs (list)", status);
 			return NULL;
@@ -480,12 +480,24 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
 
 	cl_context_properties cps[3] = { CL_CONTEXT_PLATFORM, (cl_context_properties)platform, 0 };
 
-	clState->context = clCreateContextFromType(cps, CL_DEVICE_TYPE_GPU, NULL, NULL, &status);
+	clState->context = clCreateContextFromType(cps, CL_DEVICE_TYPE_ALL, NULL, NULL, &status);
 	if (status != CL_SUCCESS) {
 		applog(LOG_ERR, "Error %d: Creating Context. (clCreateContextFromType)", status);
 		return NULL;
 	}
 
+	/////////////////////////////////////////////////////////////////
+	// Create an OpenCL command queue
+	/////////////////////////////////////////////////////////////////
+	clState->commandQueue = clCreateCommandQueue(clState->context, devices[gpu],
+						     CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &status);
+	if (status != CL_SUCCESS) /* Try again without OOE enable */
+		clState->commandQueue = clCreateCommandQueue(clState->context, devices[gpu], 0 , &status);
+	if (status != CL_SUCCESS) {
+		applog(LOG_ERR, "Error %d: Creating Command Queue. (clCreateCommandQueue)", status);
+		return NULL;
+	}
+
 	/* Check for BFI INT support. Hopefully people don't mix devices with
 	 * and without it! */
 	char * extensions = malloc(1024);
@@ -597,6 +609,8 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
 		case KL_SCRYPT:
 			strcpy(filename, SCRYPT_KERNNAME".cl");
 			strcpy(binaryfilename, SCRYPT_KERNNAME);
+			/* Scrypt only supports vector 1 */
+			gpus[gpu].vwidth = 1;
 			break;
 		case KL_NONE: /* Shouldn't happen */
 		case KL_DIABLO:
@@ -650,8 +664,8 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
 
 #ifdef USE_SCRYPT
 	if (opt_scrypt) {
-		clState->lookup_gap = 1;
-		clState->thread_concurrency = 1;
+		clState->lookup_gap = 2;
+		clState->thread_concurrency = 6144;
 	}
 #endif
 
@@ -914,25 +928,14 @@ built:
 		return NULL;
 	}
 
-	/////////////////////////////////////////////////////////////////
-	// Create an OpenCL command queue
-	/////////////////////////////////////////////////////////////////
-	clState->commandQueue = clCreateCommandQueue(clState->context, devices[gpu],
-						     CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &status);
-	if (status != CL_SUCCESS) /* Try again without OOE enable */
-		clState->commandQueue = clCreateCommandQueue(clState->context, devices[gpu], 0 , &status);
-	if (status != CL_SUCCESS) {
-		applog(LOG_ERR, "Error %d: Creating Command Queue. (clCreateCommandQueue)", status);
-		return NULL;
-	}
-
 #ifdef USE_SCRYPT
 	if (opt_scrypt) {
 		size_t ipt = (1024 / clState->lookup_gap + (1024 % clState->lookup_gap > 0));
 		size_t bufsize = 128 * ipt * clState->thread_concurrency;
 
-		clState->CLbuffer0 = clCreateBuffer(clState->context, CL_MEM_READ_ONLY, 128, NULL, &status);
+		clState->CLbuffer0 = clCreateBuffer(clState->context, CL_MEM_READ_ONLY, 80, NULL, &status);
 		clState->padbuffer8 = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, bufsize, NULL, &status);
+		clState->padbufsize = bufsize;
 	}
 #endif
 	clState->outputBuffer = clCreateBuffer(clState->context, CL_MEM_WRITE_ONLY, BUFFERSIZE, NULL, &status);

ocl.h

@@ -20,6 +20,8 @@ typedef struct {
 	cl_mem padbuffer8;
 	size_t lookup_gap;
 	size_t thread_concurrency;
+	size_t padbufsize;
+	void * cldata;
 #endif
 	bool hasBitAlign;
 	bool hasOpenCL11plus;

scrypt120713.cl

@@ -31,187 +31,187 @@ void SHA256(uint4*restrict state0,uint4*restrict state1, const uint4 block0, con
 #define G S1.z
 #define H S1.w
 
-	uint16 W;
+	uint4 W[4];
 
-	W.s0 = block0.x;
-	RND(A,B,C,D,E,F,G,H, W.s0+0x428a2f98U);
-	W.s1 = block0.y;
-	RND(H,A,B,C,D,E,F,G, W.s1+0x71374491U);
-	W.s2 = block0.z;
-	RND(G,H,A,B,C,D,E,F, W.s2+0xb5c0fbcfU);
-	W.s3 = block0.w;
-	RND(F,G,H,A,B,C,D,E, W.s3+0xe9b5dba5U);
+	W[ 0].x = block0.x;
+	RND(A,B,C,D,E,F,G,H, W[0].x+0x428a2f98U);
+	W[ 0].y = block0.y;
+	RND(H,A,B,C,D,E,F,G, W[0].y+0x71374491U);
+	W[ 0].z = block0.z;
+	RND(G,H,A,B,C,D,E,F, W[0].z+0xb5c0fbcfU);
+	W[ 0].w = block0.w;
+	RND(F,G,H,A,B,C,D,E, W[0].w+0xe9b5dba5U);
 
-	W.s4 = block1.x;
-	RND(E,F,G,H,A,B,C,D, W.s4+0x3956c25bU);
-	W.s5 = block1.y;
-	RND(D,E,F,G,H,A,B,C, W.s5+0x59f111f1U);
-	W.s6 = block1.z;
-	RND(C,D,E,F,G,H,A,B, W.s6+0x923f82a4U);
-	W.s7 = block1.w;
-	RND(B,C,D,E,F,G,H,A, W.s7+0xab1c5ed5U);
+	W[ 1].x = block1.x;
+	RND(E,F,G,H,A,B,C,D, W[1].x+0x3956c25bU);
+	W[ 1].y = block1.y;
+	RND(D,E,F,G,H,A,B,C, W[1].y+0x59f111f1U);
+	W[ 1].z = block1.z;
+	RND(C,D,E,F,G,H,A,B, W[1].z+0x923f82a4U);
+	W[ 1].w = block1.w;
+	RND(B,C,D,E,F,G,H,A, W[1].w+0xab1c5ed5U);
 
-	W.s8 = block2.x;
-	RND(A,B,C,D,E,F,G,H, W.s8+0xd807aa98U);
-	W.s9 = block2.y;
-	RND(H,A,B,C,D,E,F,G, W.s9+0x12835b01U);
-	W.sa = block2.z;
-	RND(G,H,A,B,C,D,E,F, W.sa+0x243185beU);
-	W.sb = block2.w;
-	RND(F,G,H,A,B,C,D,E, W.sb+0x550c7dc3U);
+	W[ 2].x = block2.x;
+	RND(A,B,C,D,E,F,G,H, W[2].x+0xd807aa98U);
+	W[ 2].y = block2.y;
+	RND(H,A,B,C,D,E,F,G, W[2].y+0x12835b01U);
+	W[ 2].z = block2.z;
+	RND(G,H,A,B,C,D,E,F, W[2].z+0x243185beU);
+	W[ 2].w = block2.w;
+	RND(F,G,H,A,B,C,D,E, W[2].w+0x550c7dc3U);
 
-	W.sc = block3.x;
-	RND(E,F,G,H,A,B,C,D, W.sc+0x72be5d74U);
-	W.sd = block3.y;
-	RND(D,E,F,G,H,A,B,C, W.sd+0x80deb1feU);
-	W.se = block3.z;
-	RND(C,D,E,F,G,H,A,B, W.se+0x9bdc06a7U);
-	W.sf = block3.w;
-	RND(B,C,D,E,F,G,H,A, W.sf+0xc19bf174U);
+	W[ 3].x = block3.x;
+	RND(E,F,G,H,A,B,C,D, W[3].x+0x72be5d74U);
+	W[ 3].y = block3.y;
+	RND(D,E,F,G,H,A,B,C, W[3].y+0x80deb1feU);
+	W[ 3].z = block3.z;
+	RND(C,D,E,F,G,H,A,B, W[3].z+0x9bdc06a7U);
+	W[ 3].w = block3.w;
+	RND(B,C,D,E,F,G,H,A, W[3].w+0xc19bf174U);
 
-	W.s0 += Wr1(W.se) + W.s9 + Wr2(W.s1);
-	RND(A,B,C,D,E,F,G,H, W.s0+0xe49b69c1U);
+	W[ 0].x += Wr1(W[ 3].z) + W[ 2].y + Wr2(W[ 0].y);
+	RND(A,B,C,D,E,F,G,H, W[0].x+0xe49b69c1U);
 
-	W.s1 += Wr1(W.sf) + W.sa + Wr2(W.s2);
-	RND(H,A,B,C,D,E,F,G, W.s1+0xefbe4786U);
+	W[ 0].y += Wr1(W[ 3].w) + W[ 2].z + Wr2(W[ 0].z);
+	RND(H,A,B,C,D,E,F,G, W[0].y+0xefbe4786U);
 
-	W.s2 += Wr1(W.s0) + W.sb + Wr2(W.s3);
-	RND(G,H,A,B,C,D,E,F, W.s2+0x0fc19dc6U);
+	W[ 0].z += Wr1(W[ 0].x) + W[ 2].w + Wr2(W[ 0].w);
+	RND(G,H,A,B,C,D,E,F, W[0].z+0x0fc19dc6U);
 
-	W.s3 += Wr1(W.s1) + W.sc + Wr2(W.s4);
-	RND(F,G,H,A,B,C,D,E, W.s3+0x240ca1ccU);
+	W[ 0].w += Wr1(W[ 0].y) + W[ 3].x + Wr2(W[ 1].x);
+	RND(F,G,H,A,B,C,D,E, W[0].w+0x240ca1ccU);
 
-	W.s4 += Wr1(W.s2) + W.sd + Wr2(W.s5);
-	RND(E,F,G,H,A,B,C,D, W.s4+0x2de92c6fU);
+	W[ 1].x += Wr1(W[ 0].z) + W[ 3].y + Wr2(W[ 1].y);
+	RND(E,F,G,H,A,B,C,D, W[1].x+0x2de92c6fU);
 
-	W.s5 += Wr1(W.s3) + W.se + Wr2(W.s6);
-	RND(D,E,F,G,H,A,B,C, W.s5+0x4a7484aaU);
+	W[ 1].y += Wr1(W[ 0].w) + W[ 3].z + Wr2(W[ 1].z);
+	RND(D,E,F,G,H,A,B,C, W[1].y+0x4a7484aaU);
 
-	W.s6 += Wr1(W.s4) + W.sf + Wr2(W.s7);
-	RND(C,D,E,F,G,H,A,B, W.s6+0x5cb0a9dcU);
+	W[ 1].z += Wr1(W[ 1].x) + W[ 3].w + Wr2(W[ 1].w);
+	RND(C,D,E,F,G,H,A,B, W[1].z+0x5cb0a9dcU);
 
-	W.s7 += Wr1(W.s5) + W.s0 + Wr2(W.s8);
-	RND(B,C,D,E,F,G,H,A, W.s7+0x76f988daU);
+	W[ 1].w += Wr1(W[ 1].y) + W[ 0].x + Wr2(W[ 2].x);
+	RND(B,C,D,E,F,G,H,A, W[1].w+0x76f988daU);
 
-	W.s8 += Wr1(W.s6) + W.s1 + Wr2(W.s9);
-	RND(A,B,C,D,E,F,G,H, W.s8+0x983e5152U);
+	W[ 2].x += Wr1(W[ 1].z) + W[ 0].y + Wr2(W[ 2].y);
+	RND(A,B,C,D,E,F,G,H, W[2].x+0x983e5152U);
 
-	W.s9 += Wr1(W.s7) + W.s2 + Wr2(W.sa);
-	RND(H,A,B,C,D,E,F,G, W.s9+0xa831c66dU);
+	W[ 2].y += Wr1(W[ 1].w) + W[ 0].z + Wr2(W[ 2].z);
+	RND(H,A,B,C,D,E,F,G, W[2].y+0xa831c66dU);
 
-	W.sa += Wr1(W.s8) + W.s3 + Wr2(W.sb);
-	RND(G,H,A,B,C,D,E,F, W.sa+0xb00327c8U);
+	W[ 2].z += Wr1(W[ 2].x) + W[ 0].w + Wr2(W[ 2].w);
+	RND(G,H,A,B,C,D,E,F, W[2].z+0xb00327c8U);
 
-	W.sb += Wr1(W.s9) + W.s4 + Wr2(W.sc);
-	RND(F,G,H,A,B,C,D,E, W.sb+0xbf597fc7U);
+	W[ 2].w += Wr1(W[ 2].y) + W[ 1].x + Wr2(W[ 3].x);
+	RND(F,G,H,A,B,C,D,E, W[2].w+0xbf597fc7U);
 
-	W.sc += Wr1(W.sa) + W.s5 + Wr2(W.sd);
-	RND(E,F,G,H,A,B,C,D, W.sc+0xc6e00bf3U);
+	W[ 3].x += Wr1(W[ 2].z) + W[ 1].y + Wr2(W[ 3].y);
+	RND(E,F,G,H,A,B,C,D, W[3].x+0xc6e00bf3U);
 
-	W.sd += Wr1(W.sb) + W.s6 + Wr2(W.se);
-	RND(D,E,F,G,H,A,B,C, W.sd+0xd5a79147U);
+	W[ 3].y += Wr1(W[ 2].w) + W[ 1].z + Wr2(W[ 3].z);
+	RND(D,E,F,G,H,A,B,C, W[3].y+0xd5a79147U);
 
-	W.se += Wr1(W.sc) + W.s7 + Wr2(W.sf);
-	RND(C,D,E,F,G,H,A,B, W.se+0x06ca6351U);
+	W[ 3].z += Wr1(W[ 3].x) + W[ 1].w + Wr2(W[ 3].w);
+	RND(C,D,E,F,G,H,A,B, W[3].z+0x06ca6351U);
 
-	W.sf += Wr1(W.sd) + W.s8 + Wr2(W.s0);
-	RND(B,C,D,E,F,G,H,A, W.sf+0x14292967U);
+	W[ 3].w += Wr1(W[ 3].y) + W[ 2].x + Wr2(W[ 0].x);
+	RND(B,C,D,E,F,G,H,A, W[3].w+0x14292967U);
 
-	W.s0 += Wr1(W.se) + W.s9 + Wr2(W.s1);
-	RND(A,B,C,D,E,F,G,H, W.s0+0x27b70a85U);
+	W[ 0].x += Wr1(W[ 3].z) + W[ 2].y + Wr2(W[ 0].y);
+	RND(A,B,C,D,E,F,G,H, W[0].x+0x27b70a85U);
 
-	W.s1 += Wr1(W.sf) + W.sa + Wr2(W.s2);
-	RND(H,A,B,C,D,E,F,G, W.s1+0x2e1b2138U);
+	W[ 0].y += Wr1(W[ 3].w) + W[ 2].z + Wr2(W[ 0].z);
+	RND(H,A,B,C,D,E,F,G, W[0].y+0x2e1b2138U);
 
-	W.s2 += Wr1(W.s0) + W.sb + Wr2(W.s3);
-	RND(G,H,A,B,C,D,E,F, W.s2+0x4d2c6dfcU);
+	W[ 0].z += Wr1(W[ 0].x) + W[ 2].w + Wr2(W[ 0].w);
+	RND(G,H,A,B,C,D,E,F, W[0].z+0x4d2c6dfcU);
 
-	W.s3 += Wr1(W.s1) + W.sc + Wr2(W.s4);
-	RND(F,G,H,A,B,C,D,E, W.s3+0x53380d13U);
+	W[ 0].w += Wr1(W[ 0].y) + W[ 3].x + Wr2(W[ 1].x);
+	RND(F,G,H,A,B,C,D,E, W[0].w+0x53380d13U);
 
-	W.s4 += Wr1(W.s2) + W.sd + Wr2(W.s5);
-	RND(E,F,G,H,A,B,C,D, W.s4+0x650a7354U);
+	W[ 1].x += Wr1(W[ 0].z) + W[ 3].y + Wr2(W[ 1].y);
+	RND(E,F,G,H,A,B,C,D, W[1].x+0x650a7354U);
 
-	W.s5 += Wr1(W.s3) + W.se + Wr2(W.s6);
-	RND(D,E,F,G,H,A,B,C, W.s5+0x766a0abbU);
+	W[ 1].y += Wr1(W[ 0].w) + W[ 3].z + Wr2(W[ 1].z);
+	RND(D,E,F,G,H,A,B,C, W[1].y+0x766a0abbU);
 
-	W.s6 += Wr1(W.s4) + W.sf + Wr2(W.s7);
-	RND(C,D,E,F,G,H,A,B, W.s6+0x81c2c92eU);
+	W[ 1].z += Wr1(W[ 1].x) + W[ 3].w + Wr2(W[ 1].w);
+	RND(C,D,E,F,G,H,A,B, W[1].z+0x81c2c92eU);
 
-	W.s7 += Wr1(W.s5) + W.s0 + Wr2(W.s8);
-	RND(B,C,D,E,F,G,H,A, W.s7+0x92722c85U);
+	W[ 1].w += Wr1(W[ 1].y) + W[ 0].x + Wr2(W[ 2].x);
+	RND(B,C,D,E,F,G,H,A, W[1].w+0x92722c85U);
 
-	W.s8 += Wr1(W.s6) + W.s1 + Wr2(W.s9);
-	RND(A,B,C,D,E,F,G,H, W.s8+0xa2bfe8a1U);
+	W[ 2].x += Wr1(W[ 1].z) + W[ 0].y + Wr2(W[ 2].y);
+	RND(A,B,C,D,E,F,G,H, W[2].x+0xa2bfe8a1U);
 
-	W.s9 += Wr1(W.s7) + W.s2 + Wr2(W.sa);
-	RND(H,A,B,C,D,E,F,G, W.s9+0xa81a664bU);
+	W[ 2].y += Wr1(W[ 1].w) + W[ 0].z + Wr2(W[ 2].z);
+	RND(H,A,B,C,D,E,F,G, W[2].y+0xa81a664bU);
 
-	W.sa += Wr1(W.s8) + W.s3 + Wr2(W.sb);
-	RND(G,H,A,B,C,D,E,F, W.sa+0xc24b8b70U);
+	W[ 2].z += Wr1(W[ 2].x) + W[ 0].w + Wr2(W[ 2].w);
+	RND(G,H,A,B,C,D,E,F, W[2].z+0xc24b8b70U);
 
-	W.sb += Wr1(W.s9) + W.s4 + Wr2(W.sc);
-	RND(F,G,H,A,B,C,D,E, W.sb+0xc76c51a3U);
+	W[ 2].w += Wr1(W[ 2].y) + W[ 1].x + Wr2(W[ 3].x);
+	RND(F,G,H,A,B,C,D,E, W[2].w+0xc76c51a3U);
 
-	W.sc += Wr1(W.sa) + W.s5 + Wr2(W.sd);
-	RND(E,F,G,H,A,B,C,D, W.sc+0xd192e819U);
+	W[ 3].x += Wr1(W[ 2].z) + W[ 1].y + Wr2(W[ 3].y);
+	RND(E,F,G,H,A,B,C,D, W[3].x+0xd192e819U);
 
-	W.sd += Wr1(W.sb) + W.s6 + Wr2(W.se);
-	RND(D,E,F,G,H,A,B,C, W.sd+0xd6990624U);
+	W[ 3].y += Wr1(W[ 2].w) + W[ 1].z + Wr2(W[ 3].z);
+	RND(D,E,F,G,H,A,B,C, W[3].y+0xd6990624U);
 
-	W.se += Wr1(W.sc) + W.s7 + Wr2(W.sf);
-	RND(C,D,E,F,G,H,A,B, W.se+0xf40e3585U);
+	W[ 3].z += Wr1(W[ 3].x) + W[ 1].w + Wr2(W[ 3].w);
+	RND(C,D,E,F,G,H,A,B, W[3].z+0xf40e3585U);
 
-	W.sf += Wr1(W.sd) + W.s8 + Wr2(W.s0);
-	RND(B,C,D,E,F,G,H,A, W.sf+0x106aa070U);
+	W[ 3].w += Wr1(W[ 3].y) + W[ 2].x + Wr2(W[ 0].x);
+	RND(B,C,D,E,F,G,H,A, W[3].w+0x106aa070U);
 
-	W.s0 += Wr1(W.se) + W.s9 + Wr2(W.s1);
-	RND(A,B,C,D,E,F,G,H, W.s0+0x19a4c116U);
+	W[ 0].x += Wr1(W[ 3].z) + W[ 2].y + Wr2(W[ 0].y);
+	RND(A,B,C,D,E,F,G,H, W[0].x+0x19a4c116U);
 
-	W.s1 += Wr1(W.sf) + W.sa + Wr2(W.s2);
-	RND(H,A,B,C,D,E,F,G, W.s1+0x1e376c08U);
+	W[ 0].y += Wr1(W[ 3].w) + W[ 2].z + Wr2(W[ 0].z);
+	RND(H,A,B,C,D,E,F,G, W[0].y+0x1e376c08U);
 
-	W.s2 += Wr1(W.s0) + W.sb + Wr2(W.s3);
-	RND(G,H,A,B,C,D,E,F, W.s2+0x2748774cU);
+	W[ 0].z += Wr1(W[ 0].x) + W[ 2].w + Wr2(W[ 0].w);
+	RND(G,H,A,B,C,D,E,F, W[0].z+0x2748774cU);
 
-	W.s3 += Wr1(W.s1) + W.sc + Wr2(W.s4);
-	RND(F,G,H,A,B,C,D,E, W.s3+0x34b0bcb5U);
+	W[ 0].w += Wr1(W[ 0].y) + W[ 3].x + Wr2(W[ 1].x);
+	RND(F,G,H,A,B,C,D,E, W[0].w+0x34b0bcb5U);
 
-	W.s4 += Wr1(W.s2) + W.sd + Wr2(W.s5);
-	RND(E,F,G,H,A,B,C,D, W.s4+0x391c0cb3U);
+	W[ 1].x += Wr1(W[ 0].z) + W[ 3].y + Wr2(W[ 1].y);
+	RND(E,F,G,H,A,B,C,D, W[1].x+0x391c0cb3U);
 
-	W.s5 += Wr1(W.s3) + W.se + Wr2(W.s6);
-	RND(D,E,F,G,H,A,B,C, W.s5+0x4ed8aa4aU);
+	W[ 1].y += Wr1(W[ 0].w) + W[ 3].z + Wr2(W[ 1].z);
+	RND(D,E,F,G,H,A,B,C, W[1].y+0x4ed8aa4aU);
 
-	W.s6 += Wr1(W.s4) + W.sf + Wr2(W.s7);
-	RND(C,D,E,F,G,H,A,B, W.s6+0x5b9cca4fU);
+	W[ 1].z += Wr1(W[ 1].x) + W[ 3].w + Wr2(W[ 1].w);
+	RND(C,D,E,F,G,H,A,B, W[1].z+0x5b9cca4fU);
 
-	W.s7 += Wr1(W.s5) + W.s0 + Wr2(W.s8);
-	RND(B,C,D,E,F,G,H,A, W.s7+0x682e6ff3U);
+	W[ 1].w += Wr1(W[ 1].y) + W[ 0].x + Wr2(W[ 2].x);
+	RND(B,C,D,E,F,G,H,A, W[1].w+0x682e6ff3U);
 
-	W.s8 += Wr1(W.s6) + W.s1 + Wr2(W.s9);
-	RND(A,B,C,D,E,F,G,H, W.s8+0x748f82eeU);
+	W[ 2].x += Wr1(W[ 1].z) + W[ 0].y + Wr2(W[ 2].y);
+	RND(A,B,C,D,E,F,G,H, W[2].x+0x748f82eeU);
 
-	W.s9 += Wr1(W.s7) + W.s2 + Wr2(W.sa);
-	RND(H,A,B,C,D,E,F,G, W.s9+0x78a5636fU);
+	W[ 2].y += Wr1(W[ 1].w) + W[ 0].z + Wr2(W[ 2].z);
+	RND(H,A,B,C,D,E,F,G, W[2].y+0x78a5636fU);
 
-	W.sa += Wr1(W.s8) + W.s3 + Wr2(W.sb);
-	RND(G,H,A,B,C,D,E,F, W.sa+0x84c87814U);
+	W[ 2].z += Wr1(W[ 2].x) + W[ 0].w + Wr2(W[ 2].w);
+	RND(G,H,A,B,C,D,E,F, W[2].z+0x84c87814U);
 
-	W.sb += Wr1(W.s9) + W.s4 + Wr2(W.sc);
-	RND(F,G,H,A,B,C,D,E, W.sb+0x8cc70208U);
+	W[ 2].w += Wr1(W[ 2].y) + W[ 1].x + Wr2(W[ 3].x);
+	RND(F,G,H,A,B,C,D,E, W[2].w+0x8cc70208U);
 
-	W.sc += Wr1(W.sa) + W.s5 + Wr2(W.sd);
-	RND(E,F,G,H,A,B,C,D, W.sc+0x90befffaU);
+	W[ 3].x += Wr1(W[ 2].z) + W[ 1].y + Wr2(W[ 3].y);
+	RND(E,F,G,H,A,B,C,D, W[3].x+0x90befffaU);
 
-	W.sd += Wr1(W.sb) + W.s6 + Wr2(W.se);
-	RND(D,E,F,G,H,A,B,C, W.sd+0xa4506cebU);
+	W[ 3].y += Wr1(W[ 2].w) + W[ 1].z + Wr2(W[ 3].z);
+	RND(D,E,F,G,H,A,B,C, W[3].y+0xa4506cebU);
 
-	W.se += Wr1(W.sc) + W.s7 + Wr2(W.sf);
-	RND(C,D,E,F,G,H,A,B, W.se+0xbef9a3f7U);
+	W[ 3].z += Wr1(W[ 3].x) + W[ 1].w + Wr2(W[ 3].w);
+	RND(C,D,E,F,G,H,A,B, W[3].z+0xbef9a3f7U);
 
-	W.sf += Wr1(W.sd) + W.s8 + Wr2(W.s0);
-	RND(B,C,D,E,F,G,H,A, W.sf+0xc67178f2U);
+	W[ 3].w += Wr1(W[ 3].y) + W[ 2].x + Wr2(W[ 0].x);
+	RND(B,C,D,E,F,G,H,A, W[3].w+0xc67178f2U);
 	
 #undef A
 #undef B
@@ -237,194 +237,194 @@ void SHA256_fresh(uint4*restrict state0,uint4*restrict state1, const uint4 block
 #define G (*state1).z
 #define H (*state1).w
 
-	uint16 W;
+	uint4 W[4];
 
-	W.s0 = block0.x;
-	D=0x98c7e2a2U+W.s0;
-	H=0xfc08884dU+W.s0;
+	W[0].x = block0.x;
+	D=0x98c7e2a2U+W[0].x;
+	H=0xfc08884dU+W[0].x;
 
-	W.s1 = block0.y;
-	C=0xcd2a11aeU+Tr1(D)+Ch(D,0x510e527fU,0x9b05688cU)+W.s1;
+	W[0].y = block0.y;
+	C=0xcd2a11aeU+Tr1(D)+Ch(D,0x510e527fU,0x9b05688cU)+W[0].y;
 	G=0xC3910C8EU+C+Tr2(H)+Ch(H,0xfb6feee7U,0x2a01a605U);
 
-	W.s2 = block0.z;
-	B=0x0c2e12e0U+Tr1(C)+Ch(C,D,0x510e527fU)+W.s2;
+	W[0].z = block0.z;
+	B=0x0c2e12e0U+Tr1(C)+Ch(C,D,0x510e527fU)+W[0].z;
 	F=0x4498517BU+B+Tr2(G)+Maj(G,H,0x6a09e667U);
 
-	W.s3 = block0.w;
-	A=0xa4ce148bU+Tr1(B)+Ch(B,C,D)+W.s3;
+	W[0].w = block0.w;
+	A=0xa4ce148bU+Tr1(B)+Ch(B,C,D)+W[0].w; 
 	E=0x95F61999U+A+Tr2(F)+Maj(F,G,H);
 
-	W.s4 = block1.x;
-	RND(E,F,G,H,A,B,C,D, W.s4+0x3956c25bU);
-	W.s5 = block1.y;
-	RND(D,E,F,G,H,A,B,C, W.s5+0x59f111f1U);
-	W.s6 = block1.z;
-	RND(C,D,E,F,G,H,A,B, W.s6+0x923f82a4U);
-	W.s7 = block1.w;
-	RND(B,C,D,E,F,G,H,A, W.s7+0xab1c5ed5U);
+	W[1].x = block1.x;
+	RND(E,F,G,H,A,B,C,D, W[1].x+0x3956c25bU);
+	W[1].y = block1.y;
+	RND(D,E,F,G,H,A,B,C, W[1].y+0x59f111f1U);
+	W[1].z = block1.z;
+	RND(C,D,E,F,G,H,A,B, W[1].z+0x923f82a4U);
+	W[1].w = block1.w;
+	RND(B,C,D,E,F,G,H,A, W[1].w+0xab1c5ed5U);
 	
-	W.s8 = block2.x;
-	RND(A,B,C,D,E,F,G,H, W.s8+0xd807aa98U);
-	W.s9 = block2.y;
-	RND(H,A,B,C,D,E,F,G, W.s9+0x12835b01U);
-	W.sa = block2.z;
-	RND(G,H,A,B,C,D,E,F, W.sa+0x243185beU);
-	W.sb = block2.w;
-	RND(F,G,H,A,B,C,D,E, W.sb+0x550c7dc3U);
+	W[2].x = block2.x;
+	RND(A,B,C,D,E,F,G,H, W[2].x+0xd807aa98U);
+	W[2].y = block2.y;
+	RND(H,A,B,C,D,E,F,G, W[2].y+0x12835b01U);
+	W[2].z = block2.z;
+	RND(G,H,A,B,C,D,E,F, W[2].z+0x243185beU);
+	W[2].w = block2.w;
+	RND(F,G,H,A,B,C,D,E, W[2].w+0x550c7dc3U);
 	
-	W.sc = block3.x;
-	RND(E,F,G,H,A,B,C,D, W.sc+0x72be5d74U);
-	W.sd = block3.y;
-	RND(D,E,F,G,H,A,B,C, W.sd+0x80deb1feU);
-	W.se = block3.z;
-	RND(C,D,E,F,G,H,A,B, W.se+0x9bdc06a7U);
-	W.sf = block3.w;
-	RND(B,C,D,E,F,G,H,A, W.sf+0xc19bf174U);
+	W[3].x = block3.x;
+	RND(E,F,G,H,A,B,C,D, W[3].x+0x72be5d74U);
+	W[3].y = block3.y;
+	RND(D,E,F,G,H,A,B,C, W[3].y+0x80deb1feU);
+	W[3].z = block3.z;
+	RND(C,D,E,F,G,H,A,B, W[3].z+0x9bdc06a7U);
+	W[3].w = block3.w;
+	RND(B,C,D,E,F,G,H,A, W[3].w+0xc19bf174U);
 
-	W.s0 += Wr1(W.se) + W.s9 + Wr2(W.s1);
-	RND(A,B,C,D,E,F,G,H, W.s0+0xe49b69c1U);
+	W[0].x += Wr1(W[3].z) + W[2].y + Wr2(W[0].y);
+	RND(A,B,C,D,E,F,G,H, W[0].x+0xe49b69c1U);
 
-	W.s1 += Wr1(W.sf) + W.sa + Wr2(W.s2);
-	RND(H,A,B,C,D,E,F,G, W.s1+0xefbe4786U);
+	W[0].y += Wr1(W[3].w) + W[2].z + Wr2(W[0].z);
+	RND(H,A,B,C,D,E,F,G, W[0].y+0xefbe4786U);
 
-	W.s2 += Wr1(W.s0) + W.sb + Wr2(W.s3);
-	RND(G,H,A,B,C,D,E,F, W.s2+0x0fc19dc6U);
+	W[0].z += Wr1(W[0].x) + W[2].w + Wr2(W[0].w);
+	RND(G,H,A,B,C,D,E,F, W[0].z+0x0fc19dc6U);
 
-	W.s3 += Wr1(W.s1) + W.sc + Wr2(W.s4);
-	RND(F,G,H,A,B,C,D,E, W.s3+0x240ca1ccU);
+	W[0].w += Wr1(W[0].y) + W[3].x + Wr2(W[1].x);
+	RND(F,G,H,A,B,C,D,E, W[0].w+0x240ca1ccU);
 
-	W.s4 += Wr1(W.s2) + W.sd + Wr2(W.s5);
-	RND(E,F,G,H,A,B,C,D, W.s4+0x2de92c6fU);
+	W[1].x += Wr1(W[0].z) + W[3].y + Wr2(W[1].y);
+	RND(E,F,G,H,A,B,C,D, W[1].x+0x2de92c6fU);
 
-	W.s5 += Wr1(W.s3) + W.se + Wr2(W.s6);
-	RND(D,E,F,G,H,A,B,C, W.s5+0x4a7484aaU);
+	W[1].y += Wr1(W[0].w) + W[3].z + Wr2(W[1].z);
+	RND(D,E,F,G,H,A,B,C, W[1].y+0x4a7484aaU);
 
-	W.s6 += Wr1(W.s4) + W.sf + Wr2(W.s7);
-	RND(C,D,E,F,G,H,A,B, W.s6+0x5cb0a9dcU);
+	W[1].z += Wr1(W[1].x) + W[3].w + Wr2(W[1].w);
+	RND(C,D,E,F,G,H,A,B, W[1].z+0x5cb0a9dcU);
 
-	W.s7 += Wr1(W.s5) + W.s0 + Wr2(W.s8);
-	RND(B,C,D,E,F,G,H,A, W.s7+0x76f988daU);
+	W[1].w += Wr1(W[1].y) + W[0].x + Wr2(W[2].x);
+	RND(B,C,D,E,F,G,H,A, W[1].w+0x76f988daU);
 
-	W.s8 += Wr1(W.s6) + W.s1 + Wr2(W.s9);
-	RND(A,B,C,D,E,F,G,H, W.s8+0x983e5152U);
+	W[2].x += Wr1(W[1].z) + W[0].y + Wr2(W[2].y);
+	RND(A,B,C,D,E,F,G,H, W[2].x+0x983e5152U);
 
-	W.s9 += Wr1(W.s7) + W.s2 + Wr2(W.sa);
-	RND(H,A,B,C,D,E,F,G, W.s9+0xa831c66dU);
+	W[2].y += Wr1(W[1].w) + W[0].z + Wr2(W[2].z);
+	RND(H,A,B,C,D,E,F,G, W[2].y+0xa831c66dU);
 
-	W.sa += Wr1(W.s8) + W.s3 + Wr2(W.sb);
-	RND(G,H,A,B,C,D,E,F, W.sa+0xb00327c8U);
+	W[2].z += Wr1(W[2].x) + W[0].w + Wr2(W[2].w);
+	RND(G,H,A,B,C,D,E,F, W[2].z+0xb00327c8U);
 
-	W.sb += Wr1(W.s9) + W.s4 + Wr2(W.sc);
-	RND(F,G,H,A,B,C,D,E, W.sb+0xbf597fc7U);
+	W[2].w += Wr1(W[2].y) + W[1].x + Wr2(W[3].x);
+	RND(F,G,H,A,B,C,D,E, W[2].w+0xbf597fc7U);
 
-	W.sc += Wr1(W.sa) + W.s5 + Wr2(W.sd);
-	RND(E,F,G,H,A,B,C,D, W.sc+0xc6e00bf3U);
+	W[3].x += Wr1(W[2].z) + W[1].y + Wr2(W[3].y);
+	RND(E,F,G,H,A,B,C,D, W[3].x+0xc6e00bf3U);
 
-	W.sd += Wr1(W.sb) + W.s6 + Wr2(W.se);
-	RND(D,E,F,G,H,A,B,C, W.sd+0xd5a79147U);
+	W[3].y += Wr1(W[2].w) + W[1].z + Wr2(W[3].z);
+	RND(D,E,F,G,H,A,B,C, W[3].y+0xd5a79147U);
 
-	W.se += Wr1(W.sc) + W.s7 + Wr2(W.sf);
-	RND(C,D,E,F,G,H,A,B, W.se+0x06ca6351U);
+	W[3].z += Wr1(W[3].x) + W[1].w + Wr2(W[3].w);
+	RND(C,D,E,F,G,H,A,B, W[3].z+0x06ca6351U);
 
-	W.sf += Wr1(W.sd) + W.s8 + Wr2(W.s0);
-	RND(B,C,D,E,F,G,H,A, W.sf+0x14292967U);
+	W[3].w += Wr1(W[3].y) + W[2].x + Wr2(W[0].x);
+	RND(B,C,D,E,F,G,H,A, W[3].w+0x14292967U);
 
-	W.s0 += Wr1(W.se) + W.s9 + Wr2(W.s1);
-	RND(A,B,C,D,E,F,G,H, W.s0+0x27b70a85U);
+	W[0].x += Wr1(W[3].z) + W[2].y + Wr2(W[0].y);
+	RND(A,B,C,D,E,F,G,H, W[0].x+0x27b70a85U);
 
-	W.s1 += Wr1(W.sf) + W.sa + Wr2(W.s2);
-	RND(H,A,B,C,D,E,F,G, W.s1+0x2e1b2138U);
+	W[0].y += Wr1(W[3].w) + W[2].z + Wr2(W[0].z);
+	RND(H,A,B,C,D,E,F,G, W[0].y+0x2e1b2138U);
 
-	W.s2 += Wr1(W.s0) + W.sb + Wr2(W.s3);
-	RND(G,H,A,B,C,D,E,F, W.s2+0x4d2c6dfcU);
+	W[0].z += Wr1(W[0].x) + W[2].w + Wr2(W[0].w);
+	RND(G,H,A,B,C,D,E,F, W[0].z+0x4d2c6dfcU);
 
-	W.s3 += Wr1(W.s1) + W.sc + Wr2(W.s4);
-	RND(F,G,H,A,B,C,D,E, W.s3+0x53380d13U);
+	W[0].w += Wr1(W[0].y) + W[3].x + Wr2(W[1].x);
+	RND(F,G,H,A,B,C,D,E, W[0].w+0x53380d13U);
 
-	W.s4 += Wr1(W.s2) + W.sd + Wr2(W.s5);
-	RND(E,F,G,H,A,B,C,D, W.s4+0x650a7354U);
+	W[1].x += Wr1(W[0].z) + W[3].y + Wr2(W[1].y);
+	RND(E,F,G,H,A,B,C,D, W[1].x+0x650a7354U);
 
-	W.s5 += Wr1(W.s3) + W.se + Wr2(W.s6);
-	RND(D,E,F,G,H,A,B,C, W.s5+0x766a0abbU);
+	W[1].y += Wr1(W[0].w) + W[3].z + Wr2(W[1].z);
+	RND(D,E,F,G,H,A,B,C, W[1].y+0x766a0abbU);
 
-	W.s6 += Wr1(W.s4) + W.sf + Wr2(W.s7);
-	RND(C,D,E,F,G,H,A,B, W.s6+0x81c2c92eU);
+	W[1].z += Wr1(W[1].x) + W[3].w + Wr2(W[1].w);
+	RND(C,D,E,F,G,H,A,B, W[1].z+0x81c2c92eU);
 
-	W.s7 += Wr1(W.s5) + W.s0 + Wr2(W.s8);
-	RND(B,C,D,E,F,G,H,A, W.s7+0x92722c85U);
+	W[1].w += Wr1(W[1].y) + W[0].x + Wr2(W[2].x);
+	RND(B,C,D,E,F,G,H,A, W[1].w+0x92722c85U);
 
-	W.s8 += Wr1(W.s6) + W.s1 + Wr2(W.s9);
-	RND(A,B,C,D,E,F,G,H, W.s8+0xa2bfe8a1U);
+	W[2].x += Wr1(W[1].z) + W[0].y + Wr2(W[2].y);
+	RND(A,B,C,D,E,F,G,H, W[2].x+0xa2bfe8a1U);
 
-	W.s9 += Wr1(W.s7) + W.s2 + Wr2(W.sa);
-	RND(H,A,B,C,D,E,F,G, W.s9+0xa81a664bU);
+	W[2].y += Wr1(W[1].w) + W[0].z + Wr2(W[2].z);
+	RND(H,A,B,C,D,E,F,G, W[2].y+0xa81a664bU);
 
-	W.sa += Wr1(W.s8) + W.s3 + Wr2(W.sb);
-	RND(G,H,A,B,C,D,E,F, W.sa+0xc24b8b70U);
+	W[2].z += Wr1(W[2].x) + W[0].w + Wr2(W[2].w);
+	RND(G,H,A,B,C,D,E,F, W[2].z+0xc24b8b70U);
 
-	W.sb += Wr1(W.s9) + W.s4 + Wr2(W.sc);
-	RND(F,G,H,A,B,C,D,E, W.sb+0xc76c51a3U);
+	W[2].w += Wr1(W[2].y) + W[1].x + Wr2(W[3].x);
+	RND(F,G,H,A,B,C,D,E, W[2].w+0xc76c51a3U);
 
-	W.sc += Wr1(W.sa) + W.s5 + Wr2(W.sd);
-	RND(E,F,G,H,A,B,C,D, W.sc+0xd192e819U);
+	W[3].x += Wr1(W[2].z) + W[1].y + Wr2(W[3].y);
+	RND(E,F,G,H,A,B,C,D, W[3].x+0xd192e819U);
 
-	W.sd += Wr1(W.sb) + W.s6 + Wr2(W.se);
-	RND(D,E,F,G,H,A,B,C, W.sd+0xd6990624U);
+	W[3].y += Wr1(W[2].w) + W[1].z + Wr2(W[3].z);
+	RND(D,E,F,G,H,A,B,C, W[3].y+0xd6990624U);
 
-	W.se += Wr1(W.sc) + W.s7 + Wr2(W.sf);
-	RND(C,D,E,F,G,H,A,B, W.se+0xf40e3585U);
+	W[3].z += Wr1(W[3].x) + W[1].w + Wr2(W[3].w);
+	RND(C,D,E,F,G,H,A,B, W[3].z+0xf40e3585U);
 
-	W.sf += Wr1(W.sd) + W.s8 + Wr2(W.s0);
-	RND(B,C,D,E,F,G,H,A, W.sf+0x106aa070U);
+	W[3].w += Wr1(W[3].y) + W[2].x + Wr2(W[0].x);
+	RND(B,C,D,E,F,G,H,A, W[3].w+0x106aa070U);
 
-	W.s0 += Wr1(W.se) + W.s9 + Wr2(W.s1);
-	RND(A,B,C,D,E,F,G,H, W.s0+0x19a4c116U);
+	W[0].x += Wr1(W[3].z) + W[2].y + Wr2(W[0].y);
+	RND(A,B,C,D,E,F,G,H, W[0].x+0x19a4c116U);
 
-	W.s1 += Wr1(W.sf) + W.sa + Wr2(W.s2);
-	RND(H,A,B,C,D,E,F,G, W.s1+0x1e376c08U);
+	W[0].y += Wr1(W[3].w) + W[2].z + Wr2(W[0].z);
+	RND(H,A,B,C,D,E,F,G, W[0].y+0x1e376c08U);
 
-	W.s2 += Wr1(W.s0) + W.sb + Wr2(W.s3);
-	RND(G,H,A,B,C,D,E,F, W.s2+0x2748774cU);
+	W[0].z += Wr1(W[0].x) + W[2].w + Wr2(W[0].w);
+	RND(G,H,A,B,C,D,E,F, W[0].z+0x2748774cU);
 
-	W.s3 += Wr1(W.s1) + W.sc + Wr2(W.s4);
-	RND(F,G,H,A,B,C,D,E, W.s3+0x34b0bcb5U);
+	W[0].w += Wr1(W[0].y) + W[3].x + Wr2(W[1].x);
+	RND(F,G,H,A,B,C,D,E, W[0].w+0x34b0bcb5U);
 
-	W.s4 += Wr1(W.s2) + W.sd + Wr2(W.s5);
-	RND(E,F,G,H,A,B,C,D, W.s4+0x391c0cb3U);
+	W[1].x += Wr1(W[0].z) + W[3].y + Wr2(W[1].y);
+	RND(E,F,G,H,A,B,C,D, W[1].x+0x391c0cb3U);
 
-	W.s5 += Wr1(W.s3) + W.se + Wr2(W.s6);
-	RND(D,E,F,G,H,A,B,C, W.s5+0x4ed8aa4aU);
+	W[1].y += Wr1(W[0].w) + W[3].z + Wr2(W[1].z);
+	RND(D,E,F,G,H,A,B,C, W[1].y+0x4ed8aa4aU);
 
-	W.s6 += Wr1(W.s4) + W.sf + Wr2(W.s7);
-	RND(C,D,E,F,G,H,A,B, W.s6+0x5b9cca4fU);
+	W[1].z += Wr1(W[1].x) + W[3].w + Wr2(W[1].w);
+	RND(C,D,E,F,G,H,A,B, W[1].z+0x5b9cca4fU);
 
-	W.s7 += Wr1(W.s5) + W.s0 + Wr2(W.s8);
-	RND(B,C,D,E,F,G,H,A, W.s7+0x682e6ff3U);
+	W[1].w += Wr1(W[1].y) + W[0].x + Wr2(W[2].x);
+	RND(B,C,D,E,F,G,H,A, W[1].w+0x682e6ff3U);
 
-	W.s8 += Wr1(W.s6) + W.s1 + Wr2(W.s9);
-	RND(A,B,C,D,E,F,G,H, W.s8+0x748f82eeU);
+	W[2].x += Wr1(W[1].z) + W[0].y + Wr2(W[2].y);
+	RND(A,B,C,D,E,F,G,H, W[2].x+0x748f82eeU);
 
-	W.s9 += Wr1(W.s7) + W.s2 + Wr2(W.sa);
-	RND(H,A,B,C,D,E,F,G, W.s9+0x78a5636fU);
+	W[2].y += Wr1(W[1].w) + W[0].z + Wr2(W[2].z);
+	RND(H,A,B,C,D,E,F,G, W[2].y+0x78a5636fU);
 
-	W.sa += Wr1(W.s8) + W.s3 + Wr2(W.sb);
-	RND(G,H,A,B,C,D,E,F, W.sa+0x84c87814U);
+	W[2].z += Wr1(W[2].x) + W[0].w + Wr2(W[2].w);
+	RND(G,H,A,B,C,D,E,F, W[2].z+0x84c87814U);
 
-	W.sb += Wr1(W.s9) + W.s4 + Wr2(W.sc);
-	RND(F,G,H,A,B,C,D,E, W.sb+0x8cc70208U);
+	W[2].w += Wr1(W[2].y) + W[1].x + Wr2(W[3].x);
+	RND(F,G,H,A,B,C,D,E, W[2].w+0x8cc70208U);
 
-	W.sc += Wr1(W.sa) + W.s5 + Wr2(W.sd);
-	RND(E,F,G,H,A,B,C,D, W.sc+0x90befffaU);
+	W[3].x += Wr1(W[2].z) + W[1].y + Wr2(W[3].y);
+	RND(E,F,G,H,A,B,C,D, W[3].x+0x90befffaU);
 
-	W.sd += Wr1(W.sb) + W.s6 + Wr2(W.se);
-	RND(D,E,F,G,H,A,B,C, W.sd+0xa4506cebU);
+	W[3].y += Wr1(W[2].w) + W[1].z + Wr2(W[3].z);
+	RND(D,E,F,G,H,A,B,C, W[3].y+0xa4506cebU);
 
-	W.se += Wr1(W.sc) + W.s7 + Wr2(W.sf);
-	RND(C,D,E,F,G,H,A,B, W.se+0xbef9a3f7U);
+	W[3].z += Wr1(W[3].x) + W[1].w + Wr2(W[3].w);
+	RND(C,D,E,F,G,H,A,B, W[3].z+0xbef9a3f7U);
 
-	W.sf += Wr1(W.sd) + W.s8 + Wr2(W.s0);
-	RND(B,C,D,E,F,G,H,A, W.sf+0xc67178f2U);
+	W[3].w += Wr1(W[3].y) + W[2].x + Wr2(W[0].x);
+	RND(B,C,D,E,F,G,H,A, W[3].w+0xc67178f2U);
 	
 #undef A
 #undef B
@@ -689,12 +689,13 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 #define NFLAG (0x7F)
 
 __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
-__kernel void search(__global uint4*restrict input, __global uint*restrict output, __global uint4*restrict padcache, uint4 pad0, uint4 pad1)
+__kernel void search(__global const uint4 * restrict input, __global uint*restrict output, __global uint4*restrict padcache, const uint4 midstate0, const uint4 midstate16)
 {
 	uint gid = get_global_id(0);
 	uint4 X[8];
 	uint4 tstate0, tstate1, ostate0, ostate1, tmp0, tmp1;
 	uint4 data = (uint4)(input[4].x,input[4].y,input[4].z,gid);
+	uint4 pad0 = midstate0, pad1 = midstate16;
 
 	SHA256(&pad0,&pad1, data, (uint4)(0x80000000U,0,0,0), (uint4)(0,0,0,0), (uint4)(0,0,0,0x280));
 	SHA256_fresh(&ostate0,&ostate1, pad0^0x5C5C5C5CU, pad1^0x5C5C5C5CU, 0x5C5C5C5CU, 0x5C5C5C5CU);