/**
* Author......: See docs/credits.txt
* License.....: MIT
*/

#include "inc_vendor.cl"
#include "inc_hash_constants.h"
#include "inc_hash_functions.cl"
#include "inc_types.cl"
#include "inc_common.cl"
#include "inc_hash_sha256.cl"
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"

#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"

__kernel void m13400_init (__global u32 * input, __global u32 * output, const unsigned int count)
{
const u64 gid = get_global_id (0);

//if (gid >= gid_max) return;

u32 w0[4];
u32 w1[4];
u32 w2[4];
u32 w3[4];

w0[0] = input[0+(gid*32)];
w0[1] = input[1+(gid*32)];
w0[2] = input[2+(gid*32)];
w0[3] = input[3+(gid*32)];
w1[0] = input[4+(gid*32)];
w1[1] = input[5+(gid*32)];
w1[2] = input[6+(gid*32)];
w1[3] = input[7+(gid*32)];
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
w3[3] = 0;
sha256_ctx_t ctx;


// This is just to get rid of any global memory access/initialization overhead for the benchmark

for (int i = 0; i < 1024; i += 1)
{

// Running sha256 twice
sha256_init (&ctx);
sha256_update_64 (&ctx, w0, w1, w2, w3, 0);
sha256_final (&ctx);

sha256_init (&ctx);
sha256_update_64 (&ctx, w0, w1, w2, w3, 0);
sha256_final (&ctx);


// And aes once

u32 digest[8];

__local u32 s_td0[256];
__local u32 s_td1[256];
__local u32 s_td2[256];
__local u32 s_td3[256];
__local u32 s_td4[256];

__local u32 s_te0[256];
__local u32 s_te1[256];
__local u32 s_te2[256];
__local u32 s_te3[256];
__local u32 s_te4[256];


#define KEYLEN 60

u32 ks[KEYLEN];
u32 data[4];
u32 out[4];

AES256_set_decrypt_key (ks, digest, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4);

AES256_decrypt (ks, data, out, s_td0, s_td1, s_td2, s_td3, s_td4);

}

output[0+(gid*32)] = ctx.h[0];
output[1+(gid*32)] = ctx.h[1];
output[2+(gid*32)] = ctx.h[2];
output[3+(gid*32)] = ctx.h[3];
output[4+(gid*32)] = ctx.h[4];
output[5+(gid*32)] = ctx.h[5];
output[6+(gid*32)] = ctx.h[6];
output[7+(gid*32)] = ctx.h[7];

}