#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

	uint8 ostate = {0x6a09e667U, 0xBB67AE85U, 0x3C6EF372U, 0xA54FF53AU, 0x510e527fU, 0x9b05688cU, 0x1F83D9ABU, 0x5BE0CD19U};
	uint8 tstate = {0x6a09e667U, 0xBB67AE85U, 0x3C6EF372U, 0xA54FF53AU, 0x510e527fU, 0x9b05688cU, 0x1F83D9ABU, 0x5BE0CD19U};
	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);

	Scrypt_Key(padcache, pady, input, gid, tstate, ostate);

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

	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);
}