#include "globals.h"
#include "sha.h"
#include "scrypt.h"

__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
__kernel void search(__global const uint4 *restrict input,
volatile __global uint *restrict output, __global uint4 *restrict padcache, const uint4 midstate0, const uint4 midstate16, const uint target){

	uint XA[4];
	uint XB[4];
	uint XC[4];
	uint XD[4];
	uint XE[4];
	uint XF[4];
	uint XG[4];
	uint XH[4];

	uint lnum0 = 0;
	uint lnum1 = sK[2];
	uint lnum2 = sK[4];
	uint lnum3 = sK[5];
	uint lnum4 = 0;
	bool zero = 0;
	bool one = 1;

	uint tstate0 = 0;
	uint tstate1 = 0;
	uint tstate2 = 0;
	uint tstate3 = 0;
	uint tstate4 = 0;
	uint tstate5 = 0;
	uint tstate6 = 0;
	uint tstate7 = 0;
	uint ostate0 = 0;
	uint ostate1 = 0;
	uint ostate2 = 0;
	uint ostate3 = 0;
	uint ostate4 = 0;
	uint ostate5 = 0;
	uint ostate6 = 0;
	uint ostate7 = 0;
	uint tstatebak0 = 0;
	uint tstatebak1 = 0;
	uint tstatebak2 = 0;
	uint tstatebak3 = 0;
	uint tstatebak4 = 0;
	uint tstatebak5 = 0;
	uint tstatebak6 = 0;
	uint tstatebak7 = 0;
	uint tmp00 = 0;
	uint tmp01 = 0;
	uint tmp02 = 0;
	uint tmp03 = 0;
	uint tmp04 = 0;
	uint tmp05 = 0;
	uint tmp06 = 0;
	uint tmp07 = 0;
	uint tmp08 = 0;
	uint tmp09 = 0;
	uint tmp010 = 0;
	uint tmp011 = 0;
	uint tmp012 = 0;
	uint tmp013 = 0;
	uint tmp014 = 0;
	uint tmp015 = 0;
	uint pad00 = midstate0.x;
	uint pad01 = midstate0.y;
	uint pad02 = midstate0.z;
	uint pad03 = midstate0.w;
	uint pad04 = midstate16.x;
	uint pad05 = midstate16.y;
	uint pad06 = midstate16.z;
	uint pad07 = midstate16.w;
	uint pad08 = 0;
	uint pad09 = 0;
	uint pad010 = 0;
	uint pad011 = 0;
	uint pad012 = 0;
	uint pad013 = 0;
	uint pad014 = 0;
	uint pad015 = 0;
	uint data00 = input[4].x;
	uint data01 = input[4].y;
	uint data02 = input[4].z;
	uint data03 = gid;
	uint data04 = sK[2];
	uint data05 = 0;
	uint data06 = 0;
	uint data07 = 0;
	uint data08 = 0;
	uint data09 = 0;
	uint data010 = 0;
	uint data011 = 0;
	uint data012 = 0;
	uint data013 = 0;
	uint data014 = 0;
	uint data015 = sK[3];

	SHA256(&pad00, &pad01, &pad02, &pad03, &pad04, &pad05, &pad06, &pad07, data00, data01, data02, data03,
		data04, data05, data06, data07, data08, data09, data10, data11, data12, data13, data14, data15, data16, one);

	tmp00 = pad00^sK[0]
	tmp01 = pad01^sK[0]
	tmp02 = pad02^sK[0]
	tmp03 = pad03^sK[0]
	tmp04 = pad04^sK[0]
	tmp05 = pad05^sK[0]
	tmp06 = pad06^sK[0]
	tmp07 = pad07^sK[0]
	tmp08 = sK[0];
	tmp09 = sK[0];
	tmp10 = sK[0];
	tmp11 = sK[0];
	tmp12 = sK[0];
	tmp13 = sK[0];
	tmp14 = sK[0];
	tmp15 = sK[0];

	SHA256(&0state00, &0state01, &0state02, &0state03, &0state04, &0state05, &0state06, &0state07,
		tmp00, tmp01, tmp02, tmp03, tmp04, tmp05, tmp06, tmp07, tmp08, tmp09, tmp10, tmp11, tmp12, tmp13, tmp14, tmp15, zero);

	tmp00 = pad00^sK[1]
	tmp01 = pad01^sK[1]
	tmp02 = pad02^sK[1]
	tmp03 = pad03^sK[1]
	tmp04 = pad04^sK[1]
	tmp05 = pad05^sK[1]
	tmp06 = pad06^sK[1]
	tmp07 = pad07^sK[1]
	tmp08 = sK[1];
	tmp09 = sK[1];
	tmp10 = sK[1];
	tmp11 = sK[1];
	tmp12 = sK[1];
	tmp13 = sK[1];
	tmp14 = sK[1];
	tmp15 = sK[1];

	SHA256(&tstate00, &tstate01, &tstate02, &tstate03, &tstate04, &tstate05, &tstate06, &tstate07,
		tmp00, tmp01, tmp02, tmp03, tmp04, tmp05, tmp06, tmp07, tmp08, tmp09, tmp10, tmp11, tmp12, tmp13, tmp14, tmp15, zero);

	//backup tstate
	tstatebak0 = tstate00;
	tstatebak1 = tstate01;
	tstatebak2 = tstate02;
	tstatebak3 = tstate03;
	tstatebak4 = tstate04;
	tstatebak5 = tstate05;
	tstatebak6 = tstate06;
	tstatebak7 = tstate07;

	tmp00 = input[0].x;
	tmp01 = input[0].y;
	tmp02 = input[0].z;
	tmp03 = input[0].w;
	tmp04 = input[1].x;
	tmp05 = input[1].y;
	tmp06 = input[1].z;
	tmp07 = input[1].w;
	tmp08 = input[2].x;
	tmp09 = input[2].y;
	tmp10 = input[2].z;
	tmp11 = input[2].w;
	tmp12 = input[3].x;
	tmp13 = input[3].y;
	tmp14 = input[3].z;
	tmp15 = input[3].w;


	SHA256(&tstate00, &tstate01, &tstate02, &tstate03, &tstate04, &tstate05, &tstate06, &tstate07,
		tmp00, tmp01, tmp02, tmp03, tmp04, tmp05, tmp06, tmp07, tmp08, tmp09, tmp10, tmp11, tmp12, tmp13, tmp14, tmp15, one);

	for(uint i=0; i<4; i++){
		pad00 = tstate00;
		pad01 = tstate01;
		pad02 = tstate02;
		pad03 = tstate03;
		pad04 = tstate04;
		pad05 = tstate05;
		pad06 = tstate06;
		pad07 = tstate07;
		XA[lnum0] = ostate0;
		XB[lnum0] = ostate1;
		XC[lnum0] = ostate2;
		XD[lnum0] = ostate3;
		XE[lnum0] = ostate4;
		XF[lnum0] = ostate5;
		XG[lnum0] = ostate6;
		XH[lnum0] = ostate7;

		lnum0++;
		SHA256(&pad00, &pad01, &pad02, &pad03, &pad04, &pad05, &pad06, &pad07, data00, data01, data02, data03,
			lnum0, lnum1, zero, zero, zero, zero, zero, zero, zero, zero, zero, lnum2, one);

		SHA256(&XA[i], &XB[i], &XC[i], &XD[i], &XE[i], &XF[i], &XG[i], &XH[i],
			pad00, pad01, pad02, pad03, pad04, pad05, pad06, pad07, lnum1, zero, zero, zero, zero, zero, zero, lnum3, one);
	}


	scrypt_core(XA, XB, XC, XD, XE, XF, XG, XH, padcache);

	SHA256(&tstatebak0, &tstatebak1, &tstatebak2, &tstatebak3, &tstatebak4, &tstatebak5, &tstatebak6, &tstatebak7,
		XA[0], XB[0], XC[0], XD[0], XE[0], XF[0], XG[0], XH[0], XA[1], XB[1], XC[1], XD[1], XE[1], XF[1], XG[1], XH[1], one); 

	SHA256(&tstatebak0, &tstatebak1, &tstatebak2, &tstatebak3, &tstatebak4, &tstatebak5, &tstatebak6, &tstatebak7,
		XA[2], XB[2], XC[2], XD[2], XE[2], XF[2], XG[2], XH[2], XA[3], XB[3], XC[3], XD[3], XE[3], XF[3], XG[3], XH[3], one);

	tstate00 = tstatebak0;
	tstate01 = tstatebak1;
	tstate02 = tstatebak2;
	tstate03 = tstatebak3;
	tstate04 = tstatebak4;
	tstate05 = tstatebak5;
	tstate06 = tstatebak6;
	tstate07 = tstatebak7;

#define A tstate00
#define B tstate01
#define C tstate02
#define D tstate03
#define E tstate04
#define F tstate05
#define G tstate06
#define H tstate07

	for(uint i=0; i<8; i++){
		RND(A,B,C,D,E,F,G,H, fixedWa[i]);
		RND(H,A,B,C,D,E,F,G, fixedWb[i]);
		RND(G,H,A,B,C,D,E,F, fixedWc[i]);
		RND(F,G,H,A,B,C,D,E, fixedWd[i]);
		RND(E,F,G,H,A,B,C,D, fixedWe[i]);
		RND(D,E,F,G,H,A,B,C, fixedWf[i]);
		RND(C,D,E,F,G,H,A,B, fixedWg[i]);
		RND(B,C,D,E,F,G,H,A, fixedWh[i]);
	}

#undef A
#undef B
#undef C
#undef D
#undef E
#undef F
#undef G
#undef H

	tstatebak0 += tstate00;
	tstatebak1 += tstate01;
	tstatebak2 += tstate02;
	tstatebak3 += tstate03;
	tstatebak4 += tstate04;
	tstatebak5 += tstate05;
	tstatebak6 += tstate06;
	tstatebak7 += tstate07;

	SHA256(&ostate0, &ostate1, &ostate2, &ostate3, &ostate4, &ostate5, &ostate6, &ostate7, 
		tstatebak0, tstatebak1, tstatebak2, tstatebak3, tstatebak4, tstatebak5, tstatebak6, tstatebak7, lnum1, lnum5, one);

	one = (EndianSwapa(ostate07) <= target);
	if(one)
		SETFOUND(gid);
}