#include <OpenCL/opencl.h>
#include <iostream>
#include <iomanip>
#include <fstream>
#include <memory>
#include <sys/stat.h>
#include <mach/mach_time.h>


#define SEPARATOR  ("----------------------------------------------------------------------")

uint64_t
GetCurrentTime()
{
    return mach_absolute_time();
}

double SubtractTimeInSec(uint64_t endtime, uint64_t starttime)
{
	static double conversion = 0.0;
	uint64_t difference = endtime - starttime;
	if( 0 == conversion )
	{
		mach_timebase_info_data_t timebase;
		kern_return_t kError = mach_timebase_info( &timebase );
		if( kError == 0  )
			conversion = 1e-9 * (double) timebase.numer / (double) timebase.denom;
    }

	return conversion * (double) difference;
}


uint32_t GetMaxPowerOf2(uint32_t value)
{
  uint32_t  result = 1;
  while (value && ((value & 0x1) == 0))
  {
    value >>= 1;
    result <<= 1;
  }

  return result;
}


uint32_t SwapEndianess32(uint32_t value)
{
  value = ((value << 8) & 0xFF00FF00) | ((value >> 8) & 0xFF00FF);
  return (value << 16) | (value >> 16);
}

uint64_t SwapEndianess64(uint64_t value)
{
  uint64_t  result = SwapEndianess32(value & 0xffffffff);
  result = (result << 32) | SwapEndianess32(value >> 32);

  return result;
}

uint32_t ToBCD(uint32_t value)
{
  return (((value / 1000) % 10) << 12) |
         (((value / 100) % 10) << 8) |
         (((value / 10) % 10) << 4) |
         (value % 10);
}


enum GxStat
{
  HardResetGxStat = 0x06000000,
};


enum
{
  FirstNazo = 0x0209AF08,
  SecondNazo = 0x02039df9,

  BW2NazoOffset = 0x54,

  FirstNazoOffset = 0xFC,
  SecondNazoOffset = FirstNazoOffset + 0x4C,
  ButtonMask = 0x2FFF
};

void MakeMessage(uint32_t msg[], uint32_t nazo, GxStat gxStat, uint32_t vcount,
                     uint32_t vframe, uint32_t timer0,
                     uint32_t macAddressLow, uint32_t macAddressHigh,
                     uint32_t year, uint32_t month, uint32_t day, uint32_t dow,
                     uint32_t hour, uint32_t minute, uint32_t second,
                     uint32_t heldButtons)
{
  msg[0] = SwapEndianess32(FirstNazo);
  msg[1] = SwapEndianess32(SecondNazo);
  msg[2] = SwapEndianess32(nazo);
  msg[3] = msg[4] = SwapEndianess32(nazo + BW2NazoOffset);

  msg[5] = SwapEndianess32((vcount << 16) | timer0);

  msg[6] = macAddressLow & 0xffff;

  msg[7] = (((macAddressLow >> 16) & 0xff) |
             (macAddressHigh << 8)) ^
           SwapEndianess32(gxStat ^ vframe);

  msg[8] = ((ToBCD(year) & 0xff) << 24) |
           ((ToBCD(month) & 0xff) << 16) |
           ((ToBCD(day) & 0xff) << 8) | (dow & 0xff);

  msg[9] = ((((hour >= 12) ? ToBCD(hour) + 0x40 : ToBCD(hour)) & 0xff) << 24) |
           ((ToBCD(minute) & 0xff) << 16) |
           ((ToBCD(second) & 0xff) << 8);

  msg[10] = 0;
  msg[11] = 0;

  msg[12] = SwapEndianess32(heldButtons ^ ButtonMask);

  msg[13] = 0x80000000;
  msg[14] = 0x00000000;
  msg[15] = 0x000001A0; // 416
}


static char *
LoadProgramSourceFromFile(const char *filename)
{
  struct stat statbuf;
  FILE        *fh;
  char        *source;

  fh = fopen(filename, "r");
  if (fh == 0)
    return 0;

  stat(filename, &statbuf);
  source = (char *) malloc(statbuf.st_size + 1);
  fread(source, statbuf.st_size, 1, fh);
  source[statbuf.st_size] = '\0';

  return source;
}


cl_device_id      ComputeDeviceId[2];
cl_command_queue  ComputeCommands;
cl_context        ComputeContext;
cl_program        ComputeProgram;
cl_kernel         ComputeKernel;
cl_mem            base_message_buffer;
cl_mem            output_buffer;

#define TIMER0_SEARCH_SIZE 0x1000
#define VCOUNT_SEARCH_SIZE 0x100
#define VFRAME_SEARCH_SIZE 0x8

int DoOneSearch(uint32_t *msg, uint64_t preSeed, uint32_t vframeBase)
{
  int       err = 0;

  uint64_t  t0 = GetCurrentTime();


  // Fill the input buffer with the host generated message
  //
  err = clEnqueueWriteBuffer(ComputeCommands, base_message_buffer, CL_TRUE, 0,
                             sizeof(uint32_t[16]), msg, 0, NULL, NULL);
  if (err != CL_SUCCESS)
  {
      std::cout << "Error: Failed to write to source array!" << std::endl;
      return EXIT_FAILURE;
  }

  // Set the arguments to our compute kernel
  //
  err = 0;
  err  = clSetKernelArg(ComputeKernel, 0, sizeof(cl_mem), &base_message_buffer);
  err |= clSetKernelArg(ComputeKernel, 1, sizeof(uint64_t), &preSeed);
  err |= clSetKernelArg(ComputeKernel, 2, sizeof(cl_mem), &output_buffer);
  if (err != CL_SUCCESS)
  {
    std::cout << "Error: Failed to set kernel arguments! " << err << std::endl;
    exit(1);
  }

  // Get the maximum work group size for executing the kernel on the device
  //
  size_t  local[3] = { 0, 1, 1 };
  err = clGetKernelWorkGroupInfo
          (ComputeKernel, ComputeDeviceId[0], CL_KERNEL_WORK_GROUP_SIZE,
           sizeof(local[0]), &local[0], NULL);
  if (err != CL_SUCCESS)
  {
    std::cout << "Error: Failed to retrieve kernel work group info! " << err << std::endl;
    exit(1);
  }
  //std::cout << "Kernel work group size = " << local[0] << std::endl;
  local[0] = GetMaxPowerOf2(local[0]);
  //std::cout << "Work group size = " << local[0] << std::endl;

  // Execute the kernel over the entire range of our 1d input data set
  // using the maximum number of work group items for this device
  //
  size_t  global[3] = { TIMER0_SEARCH_SIZE >> 3, VCOUNT_SEARCH_SIZE, VFRAME_SEARCH_SIZE };
  //std::cout << "Global work size = " << global[0] << std::endl;
  err = clEnqueueNDRangeKernel(ComputeCommands, ComputeKernel, 3, NULL, global,
                               local, 0, NULL, NULL);
  if (err != CL_SUCCESS)
  {
    std::cout << "Error: Failed to execute kernel! " << err << std::endl;
    return EXIT_FAILURE;
  }

  // Wait for the command commands to get serviced before reading back results
  //
  err = clFinish(ComputeCommands);
  if (err != CL_SUCCESS)
  {
    std::cout << "Error: Failed to finish commands! " << err << std::endl;
    return EXIT_FAILURE;
  }

  // Read back the results from the device to verify the output
  //
  uint8_t  *result = new uint8_t[TIMER0_SEARCH_SIZE * VCOUNT_SEARCH_SIZE * VFRAME_SEARCH_SIZE];
  err = clEnqueueReadBuffer(ComputeCommands, output_buffer, CL_TRUE, 0,
                            TIMER0_SEARCH_SIZE * VCOUNT_SEARCH_SIZE * VFRAME_SEARCH_SIZE,
                            result, 0, NULL, NULL);
  if (err != CL_SUCCESS)
  {
    std::cout << "Error: Failed to read output array! " << err << std::endl;
    return EXIT_FAILURE;
  }

  uint8_t  *rptr = result;
  for (int vf = 0; vf < VFRAME_SEARCH_SIZE; ++vf)
  {
    for (int vc = 0; vc < VCOUNT_SEARCH_SIZE; ++vc)
    {
      for (int tmr0 = 0; tmr0 < TIMER0_SEARCH_SIZE; ++tmr0)
      {
        if (*rptr++ != 0)
        {
          std::cout << std::setfill('0') << std::hex;
          std::cout << "\n*****************   Match found at timer0 = " << (tmr0 + 0x1000)
                    << ", vcount = " << vc
                    << ", vframe = " << (vframeBase | vf) << std::endl;
          std::cout << std::dec;
        }
      }
    }
  }

  std::cout << std::dec;

  uint64_t  t1 = GetCurrentTime();

  double t = SubtractTimeInSec(t1, t0);
  std::cout << "Exec Time:  " << (1000.0 * t) << " ms";

  delete[] result;

  return EXIT_SUCCESS;
}


int DoFullSearch()
{
  uint32_t  year = 2012;
  uint32_t  month = 10;
  uint32_t  day = 12;
  uint32_t  dayOfWeek = 5;
  uint32_t  hour = 20;
  uint32_t  min = 0;
  uint32_t  second = 23;
  uint32_t  macHigh = 0x8C56C5;
  uint32_t  macLow = 0xE57B30;
  uint64_t  fullseed = 0x69E235BA6219A09CULL;

  uint64_t  preSeed = SwapEndianess64((fullseed * 0xDEDCEDAE9638806DULL) + 0x9B1AE6E9A384E6F9ULL);

  std::cout << std::setfill('0');
  std::cout << std::hex;
  std::cout << "Searching for pre-seed " << std::setw(16) << preSeed
            << " of actual seed " << std::setw(16) << fullseed << std::endl;

  uint32_t  msg[16];
  uint32_t  nazo = 0x027a6a90;

  for (uint32_t t = 2; t < 9; ++t)
  {
    while (nazo > 0x027A3000)
    {
      std::cout << std::hex;
      std::cout << "\nNazo = 0x" << nazo << ": Second = ";
      std::cout << std::dec << (second - t);

      std::cout << " vframeBase = 0x0, ";
      MakeMessage(msg, nazo, HardResetGxStat, 0, 0, 0, macLow, macHigh,
                  year, month, day, dayOfWeek, hour, min, second - t, 0);
      std::cout.flush();
      if (DoOneSearch(msg, preSeed, 0) != EXIT_SUCCESS)
      {
        std::cerr << "Error encountered - terminating" << std::endl;
        return EXIT_FAILURE;
      }

      std::cout << "  -  vframeBase = 0x8, ";
      std::cout.flush();
      MakeMessage(msg, nazo, HardResetGxStat, 0, 8, 0, macLow, macHigh,
                  year, month, day, dayOfWeek, hour, min, second - t, 0);
      if (DoOneSearch(msg, preSeed, 8) != EXIT_SUCCESS)
      {
        std::cerr << "Error encountered - terminating" << std::endl;
        return EXIT_FAILURE;
      }
      std::cout.flush();

      nazo -= 0x20;
    }

    nazo = 0x027A8010;
  }

  return EXIT_SUCCESS;
}


int main (int argc, char * const argv[])
{
  int       err = 0;

  // Connect to a GPU compute device
  //
  err = clGetDeviceIDs(NULL, CL_DEVICE_TYPE_CPU, 1, ComputeDeviceId, NULL);
  if (err != CL_SUCCESS)
  {
      std::cout << "Error: Failed to locate a compute device!" << std::endl;
      return EXIT_FAILURE;
  }
  //ComputeDeviceId[0] = ComputeDeviceId[1];

  size_t  returned_size;
  cl_char vendor_name[1024] = {0};
  cl_char device_name[1024] = {0};
  err = clGetDeviceInfo(ComputeDeviceId[0], CL_DEVICE_VENDOR, sizeof(vendor_name),
                        vendor_name, &returned_size);
  err|= clGetDeviceInfo(ComputeDeviceId[0], CL_DEVICE_NAME, sizeof(device_name),
                        device_name, &returned_size);
  if (err != CL_SUCCESS)
  {
      std::cout << "Error: Failed to retrieve device info!" << std::endl;
      return EXIT_FAILURE;
  }

  std::cout << SEPARATOR << std::endl;
  std::cout << "Connecting to " << vendor_name << ' ' << device_name << std::endl;

  // Load the compute program from disk into a cstring buffer
  //
  std::cout << SEPARATOR << std::endl;
  std::cout << "Loading program '../../sha1.cl' ..." << std::endl;
  std::cout << SEPARATOR << std::endl;

  char  *source = LoadProgramSourceFromFile("../../sha1.cl");

  // Create a compute ComputeContext
  //
  ComputeContext = clCreateContext(0, 1, &ComputeDeviceId[0], NULL, NULL, &err);
  if (!ComputeContext)
  {
      std::cout << "Error: Failed to create a compute ComputeContext!" << std::endl;
      return EXIT_FAILURE;
  }

  // Create a command queue
  //
  ComputeCommands = clCreateCommandQueue(ComputeContext, ComputeDeviceId[0], 0, &err);
  if (!ComputeCommands)
  {
      std::cout << "Error: Failed to create a command ComputeCommands!" << std::endl;
      return EXIT_FAILURE;
  }

  // Create the compute program from the source buffer
  //
  ComputeProgram = clCreateProgramWithSource(ComputeContext, 1,
                                             (const char **) &source,
                                             NULL, &err);
  if (!ComputeProgram || err != CL_SUCCESS)
  {
      std::cout << "Error: Failed to create compute program!" << std::endl;
      return EXIT_FAILURE;
  }

  // Build the program executable
  //
  err = clBuildProgram(ComputeProgram, 0, NULL, NULL, NULL, NULL);
  if (err != CL_SUCCESS)
  {
      size_t length;
      char build_log[2048];
      std::cout << "Error: Failed to build program executable!" << std::endl;
      clGetProgramBuildInfo(ComputeProgram, ComputeDeviceId[0], CL_PROGRAM_BUILD_LOG, sizeof(build_log), build_log, &length);
      std::cout << build_log << std::endl;
      return EXIT_FAILURE;
  }

  // Create the compute kernel from within the program
  //
  ComputeKernel = clCreateKernel(ComputeProgram, "sha1_v0", &err);
  if (!ComputeKernel || err != CL_SUCCESS)
  {
      std::cout << "Error: Failed to create compute kernel!" << std::endl;
      return EXIT_FAILURE;
  }

  // Create the input buffer on the device
  //
  size_t buffer_size = sizeof(uint32_t) * 16;
  base_message_buffer = clCreateBuffer(ComputeContext, CL_MEM_READ_ONLY, buffer_size, NULL, NULL);
  if (!base_message_buffer)
  {
      std::cout << "Error: Failed to allocate base message buffer on device!" << std::endl;
      return EXIT_FAILURE;
  }

  // Create the output buffer on the device
  //
  buffer_size = TIMER0_SEARCH_SIZE * VCOUNT_SEARCH_SIZE * VFRAME_SEARCH_SIZE;
  output_buffer = clCreateBuffer(ComputeContext, CL_MEM_WRITE_ONLY, buffer_size, NULL, &err);
  if (!output_buffer)
  {
      std::cout << "Error: Failed to allocate result buffer on device! " << err << std::endl;
      return EXIT_FAILURE;
  }

  DoFullSearch();

  // Shutdown and cleanup
  //
  clReleaseKernel(ComputeKernel);
  clReleaseProgram(ComputeProgram);
  clReleaseMemObject(base_message_buffer);
  clReleaseMemObject(output_buffer);
  clReleaseCommandQueue(ComputeCommands);
  clReleaseContext(ComputeContext);

  return 0;
}