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

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

#ifdef GOFFSET
	uint gid = get_global_id(0);
#else
	uint gid = input->f + get_global_id(0);
#endif

//	__local uint8 fresh;
//	fresh = (uint8)(0x6a09e667U, 0xBB67AE85U, 0x3C6EF372U, 0xA54FF53AU, 0x510e527fU, 0x9b05688cU, 0x1F83D9ABU, 0x5BE0CD19U);

	uint8 ostate = {0x6a09e667U, 0xBB67AE85U, 0x3C6EF372U, 0xA54FF53AU, 0x510e527fU, 0x9b05688cU, 0x1F83D9ABU, 0x5BE0CD19U};
	uint8 tstate = {0x6a09e667U, 0xBB67AE85U, 0x3C6EF372U, 0xA54FF53AU, 0x510e527fU, 0x9b05688cU, 0x1F83D9ABU, 0x5BE0CD19U};

//uint8 tstate;
//uint8 ostate, tstate;



uint8 tmpa = {input->c, input->d, input->e, gid, SK00, ZERO, ZERO, ZERO};
uint8 tmpb = {midstate0.s0, midstate0.s1, midstate0.s2, midstate0.s3, midstate16.s0, midstate16.s1, midstate16.s2, midstate16.s3};
uint8 tmpc = {ZERO, ZERO, ZERO, ZERO, ZERO, ZERO, ZERO, SK01};
	uint padx = gid%CONCURRENT_THREADS;
	uint pady = padx<<1;
	uint padz = pady+1;

    SHA256(&tmpb, &tmpa, &tmpc);

tmpa = tmpb^SK02;
tmpc = SK02;
SHA256(&ostate, &tmpa, &tmpc);
tmpa = tmpb^SK03;
tmpc = SK03;
SHA256(&tstate, &tmpa, &tmpc);


//Int_SHA(tmpb, &tstate, &ostate);
//    uint8 tmpa = midsha^SK02;
//    uint8 tmpb = SK02;
/*
	tmpa = tmpb^SK02;
	tmpc = SK02;
    for(uint i=0; i<2; i++){
        SHA256((i==0 ? &ostate : &tstate), &tmpa, &tmpc);

        if(i==1)
            break;
        tmpa = tmpb^SK03;
        tmpc = SK03;
    }
*/

//    SHA256(&tmpc, &tmpa, &tmpb);
//	tmpc = SK02;
//   	tmpa = tmpb^tmpc;
    //tmpc = SK02;
//    SHA256(&ostate, &tmpa, &tmpc);
//	tmpc = SK03;
//   tmpa = tmpb^tmpc;
    //tmpc = SK03;
//    SHA256(&tstate, &tmpa, &tmpc);

//Int_SHA((uint8)(midstate0.s0, midstate0.s1, midstate0.s2, midstate0.s3, midstate16.s0, midstate16.s1, midstate16.s2, midstate16.s3),
//(uint8)(input->c, input->d, input->e, gid, SK00, ZERO, ZERO, ZERO), &tstate, &ostate);

	//tmpc = SK02;
	//tstate = fresh;
//	for(uint i=0; i<2; i++){
//		tmpa = i==0 ? tmpb^SK02 : tmpb^SK03;
//		tmpc = i==0 ? SK02 : SK03;
/*
if(i){
tmpa = tmpb^SK03;
tmpc = SK03;
}else{
tmpa = tmpb^SK02;
tmpc = SK02;
//tstate = fresh;
}
*/

//tstate = ostate;
		//tmpa = tmpb^tmpc;

//		SHA256(&tstate, &tmpa, &tmpc);
//		if(i==1)
//			break;
//		tmpc = ostate;
//		ostate = tstate;
//		tstate = tmpc;
		//tmpc = SK03;
//	}
/*
for(uint i=0; i<3; i++){
SHA256(&tmpa, &tmpb, &tmpc);
if(i==2)
break;

if(i==0){
tmpb = tmpa^SK02;
tmpc = SK02;
}else{
tmpb = tmpa^SK03;
tmpc = SK03;
}

}
*/
	Scrypt_Key(padcache, pady, input, gid, tstate, ostate);

Shittify(&padcache[pady], &padcache[padz], 1);
	scrypt_core(padcache, padx);
Shittify(&padcache[pady], &padcache[padz], 0);

//Tstate_SHAd(&tstate, padcache, pady);

//Tstate_SHAd(&tstate, &padcache[pady]);

/*
tmpa = padcache[pady].lo;
tmpb = padcache[pady++].hi;

for(uint i=0; i<2; i++){
    SHA256(&tstate, &tmpa, &tmpb);
    if(i==1)
        break;
    tmpa = padcache[pady].lo;
    tmpb = padcache[pady].hi;
}
//SHA256(&tstate, &tmpa, &tmpb);
*/

for(uint i=0; i<2; i++){
tmpa = padcache[pady].lo;
tmpb = padcache[pady++].hi;
SHA256(&tstate, &tmpa, &tmpb);
}

Final_SHAd(&tstate, &ostate);
	if( (EndianSwapa((ostate.s7)) <= target) )
		SETFOUND(gid);
}