/*
* 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