/*
 * FugueCoin kernel implementation.
 *
 * ==========================(LICENSE BEGIN)============================
 *
 * Copyright (c) 2014  phm
 * 
 * Permission is hereby granted, free of charge, to any person obtaining
 * a copy of this software and associated documentation files (the
 * "Software"), to deal in the Software without restriction, including
 * without limitation the rights to use, copy, modify, merge, publish,
 * distribute, sublicense, and/or sell copies of the Software, and to
 * permit persons to whom the Software is furnished to do so, subject to
 * the following conditions:
 * 
 * The above copyright notice and this permission notice shall be
 * included in all copies or substantial portions of the Software.
 * 
 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
 * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
 * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
 * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
 * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
 * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
 * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
 *
 * ===========================(LICENSE END)=============================
 *
 * @author   phm <phm@inbox.com>
 */

#ifndef FUGUECOIN_CL
#define FUGUECOIN_CL

#if __ENDIAN_LITTLE__
#define SPH_LITTLE_ENDIAN 1
#else
#define SPH_BIG_ENDIAN 1
#endif

#define SPH_UPTR sph_u64

typedef unsigned int sph_u32;
typedef int sph_s32;
#ifndef __OPENCL_VERSION__
typedef unsigned long long sph_u64;
typedef long long sph_s64;
#else
typedef unsigned long sph_u64;
typedef long sph_s64;
#endif

#define SPH_64 1
#define SPH_64_TRUE 1

#define SPH_C32(x)    ((sph_u32)(x ## U))
#define SPH_T32(x)    ((x) & SPH_C32(0xFFFFFFFF))
#define SPH_ROTL32(x, n)   SPH_T32(((x) << (n)) | ((x) >> (32 - (n))))
#define SPH_ROTR32(x, n)   SPH_ROTL32(x, (32 - (n)))

#define SPH_C64(x)    ((sph_u64)(x ## UL))
#define SPH_T64(x)    ((x) & SPH_C64(0xFFFFFFFFFFFFFFFF))
#define SPH_ROTL64(x, n)   SPH_T64(((x) << (n)) | ((x) >> (64 - (n))))
#define SPH_ROTR64(x, n)   SPH_ROTL64(x, (64 - (n)))

#include "fugue.cl"

#define SWAP4(x) as_uint(as_uchar4(x).wzyx)
#define SWAP8(x) as_ulong(as_uchar8(x).s76543210)

#if SPH_BIG_ENDIAN
  #define DEC32BE(x) (*(const __global sph_u32 *) (x))
#else
  #define DEC32BE(x) SWAP4(*(const __global sph_u32 *) (x))
#endif

__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
__kernel void search(__global unsigned char* input, volatile __global uint* output, const ulong target)
{
  uint gid = get_global_id(0);

  //mixtab
  __local sph_u32 mixtab0[256], mixtab1[256], mixtab2[256], mixtab3[256];
  int init = get_local_id(0);
  int step = get_local_size(0);
  for (int i = init; i < 256; i += step)
  {
    mixtab0[i] = mixtab0_c[i];
    mixtab1[i] = mixtab1_c[i];
    mixtab2[i] = mixtab2_c[i];
    mixtab3[i] = mixtab3_c[i];
  }
  barrier(CLK_GLOBAL_MEM_FENCE);

  sph_u32 S00 = 0, S01 = 0, S02 = 0, S03 = 0, S04 = 0, S05 = 0, S06 = 0, S07 = 0, S08 = 0, S09 = 0; \
  sph_u32 S10 = 0, S11 = 0, S12 = 0, S13 = 0, S14 = 0, S15 = 0, S16 = 0, S17 = 0, S18 = 0, S19 = 0; \
  sph_u32 S20 = 0, S21 = 0, S22 = IV256[0], S23 = IV256[1], S24 = IV256[2], S25 = IV256[3], S26 = IV256[4], S27 = IV256[5], S28 = IV256[6], S29 = IV256[7];

  FUGUE256_5(DEC32BE(input + 0x0), DEC32BE(input + 0x4), DEC32BE(input + 0x8), DEC32BE(input + 0xc), DEC32BE(input + 0x10));
  FUGUE256_5(DEC32BE(input + 0x14), DEC32BE(input + 0x18), DEC32BE(input + 0x1c), DEC32BE(input + 0x20), DEC32BE(input + 0x24));
  FUGUE256_5(DEC32BE(input + 0x28), DEC32BE(input + 0x2c), DEC32BE(input + 0x30), DEC32BE(input + 0x34), DEC32BE(input + 0x38));
  FUGUE256_4(DEC32BE(input + 0x3c), DEC32BE(input + 0x40), DEC32BE(input + 0x44), DEC32BE(input + 0x48));

  TIX2(SWAP4(gid), S06, S07, S14, S16, S00);
  CMIX30(S03, S04, S05, S07, S08, S09, S18, S19, S20);
  SMIX(S03, S04, S05, S06);
  CMIX30(S00, S01, S02, S04, S05, S06, S15, S16, S17);
  SMIX(S00, S01, S02, S03);

  TIX2(0, S00, S01, S08, S10, S24);
  CMIX30(S27, S28, S29, S01, S02, S03, S12, S13, S14);
  SMIX(S27, S28, S29, S00);
  CMIX30(S24, S25, S26, S28, S29, S00, S09, S10, S11);
  SMIX(S24, S25, S26, S27);

  TIX2(0x280, S24, S25, S02, S04, S18);
  CMIX30(S21, S22, S23, S25, S26, S27, S06, S07, S08);
  SMIX(S21, S22, S23, S24);
  CMIX30(S18, S19, S20, S22, S23, S24, S03, S04, S05);
  SMIX(S18, S19, S20, S21);

  CMIX30(S15, S16, S17, S19, S20, S21, S00, S01, S02);
  SMIX(S15, S16, S17, S18);
  CMIX30(S12, S13, S14, S16, S17, S18, S27, S28, S29);
  SMIX(S12, S13, S14, S15);
  CMIX30(S09, S10, S11, S13, S14, S15, S24, S25, S26);
  SMIX(S09, S10, S11, S12);
  CMIX30(S06, S07, S08, S10, S11, S12, S21, S22, S23);
  SMIX(S06, S07, S08, S09);
  CMIX30(S03, S04, S05, S07, S08, S09, S18, S19, S20);
  SMIX(S03, S04, S05, S06);
  CMIX30(S00, S01, S02, S04, S05, S06, S15, S16, S17);
  SMIX(S00, S01, S02, S03);
  CMIX30(S27, S28, S29, S01, S02, S03, S12, S13, S14);
  SMIX(S27, S28, S29, S00);
  CMIX30(S24, S25, S26, S28, S29, S00, S09, S10, S11);
  SMIX(S24, S25, S26, S27);
  CMIX30(S21, S22, S23, S25, S26, S27, S06, S07, S08);
  SMIX(S21, S22, S23, S24);
  CMIX30(S18, S19, S20, S22, S23, S24, S03, S04, S05);
  SMIX(S18, S19, S20, S21);
  S22 ^= S18;
  S03 ^= S18;
  SMIX(S03, S04, S05, S06);
  S07 ^= S03;
  S19 ^= S03;
  SMIX(S19, S20, S21, S22);
  S23 ^= S19;
  S04 ^= S19;
  SMIX(S04, S05, S06, S07);
  S08 ^= S04;
  S20 ^= S04;
  SMIX(S20, S21, S22, S23);
  S24 ^= S20;
  S05 ^= S20;
  SMIX(S05, S06, S07, S08);
  S09 ^= S05;
  S21 ^= S05;
  SMIX(S21, S22, S23, S24);
  S25 ^= S21;
  S06 ^= S21;
  SMIX(S06, S07, S08, S09);
  S10 ^= S06;
  S22 ^= S06;
  SMIX(S22, S23, S24, S25);
  S26 ^= S22;
  S07 ^= S22;
  SMIX(S07, S08, S09, S10);
  S11 ^= S07;
  S23 ^= S07;
  SMIX(S23, S24, S25, S26);
  S27 ^= S23;
  S08 ^= S23;
  SMIX(S08, S09, S10, S11);
  S12 ^= S08;
  S24 ^= S08;
  SMIX(S24, S25, S26, S27);
  S28 ^= S24;
  S09 ^= S24;
  SMIX(S09, S10, S11, S12);
  S13 ^= S09;
  S25 ^= S09;
  SMIX(S25, S26, S27, S28);
  S29 ^= S25;
  S10 ^= S25;
  SMIX(S10, S11, S12, S13);
  S14 ^= S10;
  S26 ^= S10;
  SMIX(S26, S27, S28, S29);
  S00 ^= S26;
  S11 ^= S26;
  SMIX(S11, S12, S13, S14);
  S15 ^= S11;
  S27 ^= S11;
  SMIX(S27, S28, S29, S00);
  S01 ^= S27;
  S12 ^= S27;
  SMIX(S12, S13, S14, S15);
  S16 ^= S12;
  S28 ^= S12;
  SMIX(S28, S29, S00, S01);
  S02 ^= S28;
  S13 ^= S28;
  SMIX(S13, S14, S15, S16);
  S17 ^= S13;
  S29 ^= S13;
  SMIX(S29, S00, S01, S02);
  S03 ^= S29;
  S14 ^= S29;
  SMIX(S14, S15, S16, S17);
  S18 ^= S14;
  S00 ^= S14;
  SMIX(S00, S01, S02, S03);
  S04 ^= S00;
  S15 ^= S00;
  SMIX(S15, S16, S17, S18);
  S19 ^= S15;
  S01 ^= S15;
  SMIX(S01, S02, S03, S04);

  S05 ^= S01;
  S16 ^= S01;

  bool result = ((((sph_u64) SWAP4(S19) << 32) | SWAP4(S18)) <= target);
  if (result)
    output[output[0xFF]++] = SWAP4(gid);
}

#endif // FUGUECOIN_CL