// leftrotate function definition
#define LEFTROTATE(x, c) (((x) << (c)) | ((x) >> (32 - (c))))

#define F(x,y,z) ((x & y) | ((~x) & z))
#define G(x,y,z) ((x & z) | ((~z) & y))
#define H(x,y,z) (x ^ y ^ z)
#define I(x,y,z) (y ^ (x | (~z)))

// This version of the round shifts the interpretation of a,b,c,d by one
// and must be called with v/x/y/z in a matching shuffle pattern.
// Every four Rounds, a,b,c,d are back to their original interpretation,
// thogh, so it all works out in the end (we have 64 rounds per block).
#define ROUND_INPLACE_VIA_SHIFT(w, r, k, v, x, y, z, func)       \
{                                                                \
    v += func(x,y,z) + w + k;                                    \
    v = x + LEFTROTATE(v, r);                                    \
}

// This version ignores the mapping of a/b/c/d to v/x/y/z and simply
// uses a temporary variable to keep the interpretation of a/b/c/d
// consistent.  Whether this one or the previous one performs better
// probably depends on the compiler....
#define ROUND_USING_TEMP_VARS(w, r, k, v, x, y, z, func)         \
{                                                                \
    a = a + func(b,c,d) + k + w;                                 \
    unsigned int temp = d;                                       \
    d = c;                                                       \
    c = b;                                                       \
    b = b + LEFTROTATE(a, r);                                    \
    a = temp;                                                    \
}

// Here, we pick which style of ROUND we use.
#define ROUND ROUND_USING_TEMP_VARS
//#define ROUND ROUND_INPLACE_VIA_SHIFT

/// NOTE: this really only allows a length up to 7 bytes, not 8, because
/// we need to start the padding in the first byte following the message,
/// and we only have two words to work with here....
/// It also assumes words[] has all zero bits except the chars of interest.
void md5_2words(unsigned int *words, unsigned int len,
                       unsigned int *digest);
// ****************************************************************************
// Function:  IndexToKey
//
// Purpose:
///   For a given index in the keyspace, find the actual key string
///   which is at that index.
//
// Arguments:
//   index         index in key space
//   byteLength    number of bytes in a key
//   valsPerByte   number of values each byte can take on
//   vals          output key string
//
// Programmer:  Jeremy Meredith
// Creation:    July 23, 2014
//
// Modifications:
// ****************************************************************************
inline void IndexToKey(unsigned int index, int byteLength, int valsPerByte,
                       unsigned char vals[8])
{
    vals[0] = index % valsPerByte;
    index /= valsPerByte;

    vals[1] = index % valsPerByte;
    index /= valsPerByte;

    vals[2] = index % valsPerByte;
    index /= valsPerByte;

    vals[3] = index % valsPerByte;
    index /= valsPerByte;

    vals[4] = index % valsPerByte;
    index /= valsPerByte;

    vals[5] = index % valsPerByte;
    index /= valsPerByte;

    vals[6] = index % valsPerByte;
    index /= valsPerByte;

    vals[7] = index % valsPerByte;
    index /= valsPerByte;
}


// ****************************************************************************
// Function:  FindKeyWithDigest_Kernel
//
// Purpose:
///   Within each thread of a GPU, search part of the key space
///   to find a key with the given digest.
//
// Arguments:
//   searchDigest    the digest to search for
//   keyspace        the size of the key space to search
//   byteLength      number of bytes in a key
//   valsPerByte     number of values each byte can take on
//   foundIndex      output - the index of the found key (if found)
//   foundKey        output - the string of the found key (if found)
//   foundDigest     output - the digest of the found key (if found)
//
// Programmer:  Jeremy Meredith
// Creation:    July 23, 2014
//
// Modifications:
// ****************************************************************************
__kernel void
FindKeyWithDigest_Kernel(unsigned int searchDigest0,
                         unsigned int searchDigest1,
                         unsigned int searchDigest2,
                         unsigned int searchDigest3,
                         int keyspace,
                         int byteLength, int valsPerByte,
                         __global int *foundIndex,
                         __global unsigned char *foundKey,
                         __global unsigned int *foundDigest)
{
    int threadid = (get_group_id(0)*get_local_size(0)) + get_local_id(0);

    int startindex = threadid * valsPerByte;
    unsigned char key[8] = {0,0,0,0, 0,0,0,0};
    IndexToKey(startindex, byteLength, valsPerByte, key);

    for (int j=0; j < valsPerByte && startindex+j < keyspace; ++j)
    {
        unsigned int digest[4];
        md5_2words((unsigned int*)key, byteLength, digest);
        if (digest[0] == searchDigest0 &&
            digest[1] == searchDigest1 &&
            digest[2] == searchDigest2 &&
            digest[3] == searchDigest3)
        {
            *foundIndex = startindex + j;
            foundKey[0] = key[0];
            foundKey[1] = key[1];
            foundKey[2] = key[2];
            foundKey[3] = key[3];
            foundKey[4] = key[4];
            foundKey[5] = key[5];
            foundKey[6] = key[6];
            foundKey[7] = key[7];
            foundDigest[0] = digest[0];
            foundDigest[1] = digest[1];
            foundDigest[2] = digest[2];
            foundDigest[3] = digest[3];
        }
        ++key[0];
    }   
}


