SHARE
TWEET

Untitled

a guest Jan 20th, 2019 74 Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
  1.  
  2. #include <OpenCL/opencl.h>
  3. #include <iostream>
  4. #include <iomanip>
  5. #include <fstream>
  6. #include <memory>
  7. #include <sys/stat.h>
  8. #include <mach/mach_time.h>
  9.  
  10.  
  11. #define SEPARATOR  ("----------------------------------------------------------------------")
  12.  
  13. uint64_t
  14. GetCurrentTime()
  15. {
  16.     return mach_absolute_time();
  17. }
  18.    
  19. double SubtractTimeInSec(uint64_t endtime, uint64_t starttime)
  20. {    
  21.     static double conversion = 0.0;
  22.     uint64_t difference = endtime - starttime;
  23.     if( 0 == conversion )
  24.     {
  25.         mach_timebase_info_data_t timebase;
  26.         kern_return_t kError = mach_timebase_info( &timebase );
  27.         if( kError == 0  )
  28.             conversion = 1e-9 * (double) timebase.numer / (double) timebase.denom;
  29.     }
  30.        
  31.     return conversion * (double) difference;
  32. }
  33.  
  34.  
  35. uint32_t GetMaxPowerOf2(uint32_t value)
  36. {
  37.   uint32_t  result = 1;
  38.   while (value && ((value & 0x1) == 0))
  39.   {
  40.     value >>= 1;
  41.     result <<= 1;
  42.   }
  43.  
  44.   return result;
  45. }
  46.  
  47.  
  48. uint32_t SwapEndianess32(uint32_t value)
  49. {
  50.   value = ((value << 8) & 0xFF00FF00) | ((value >> 8) & 0xFF00FF);
  51.   return (value << 16) | (value >> 16);
  52. }
  53.  
  54. uint64_t SwapEndianess64(uint64_t value)
  55. {
  56.   uint64_t  result = SwapEndianess32(value & 0xffffffff);
  57.   result = (result << 32) | SwapEndianess32(value >> 32);
  58.  
  59.   return result;
  60. }
  61.  
  62. uint32_t ToBCD(uint32_t value)
  63. {
  64.   return (((value / 1000) % 10) << 12) |
  65.          (((value / 100) % 10) << 8) |
  66.          (((value / 10) % 10) << 4) |
  67.          (value % 10);
  68. }
  69.  
  70.  
  71. enum GxStat
  72. {
  73.   HardResetGxStat = 0x06000000,
  74. };
  75.  
  76.  
  77. enum
  78. {
  79.   FirstNazo = 0x0209AF08,
  80.   SecondNazo = 0x02039df9,
  81.  
  82.   BW2NazoOffset = 0x54,
  83.  
  84.   FirstNazoOffset = 0xFC,
  85.   SecondNazoOffset = FirstNazoOffset + 0x4C,
  86.   ButtonMask = 0x2FFF
  87. };
  88.  
  89. void MakeMessage(uint32_t msg[], uint32_t nazo, GxStat gxStat, uint32_t vcount,
  90.                      uint32_t vframe, uint32_t timer0,
  91.                      uint32_t macAddressLow, uint32_t macAddressHigh,
  92.                      uint32_t year, uint32_t month, uint32_t day, uint32_t dow,
  93.                      uint32_t hour, uint32_t minute, uint32_t second,
  94.                      uint32_t heldButtons)
  95. {
  96.   msg[0] = SwapEndianess32(FirstNazo);
  97.   msg[1] = SwapEndianess32(SecondNazo);
  98.   msg[2] = SwapEndianess32(nazo);
  99.   msg[3] = msg[4] = SwapEndianess32(nazo + BW2NazoOffset);
  100.  
  101.   msg[5] = SwapEndianess32((vcount << 16) | timer0);
  102.  
  103.   msg[6] = macAddressLow & 0xffff;
  104.  
  105.   msg[7] = (((macAddressLow >> 16) & 0xff) |
  106.              (macAddressHigh << 8)) ^
  107.            SwapEndianess32(gxStat ^ vframe);
  108.  
  109.   msg[8] = ((ToBCD(year) & 0xff) << 24) |
  110.            ((ToBCD(month) & 0xff) << 16) |
  111.            ((ToBCD(day) & 0xff) << 8) | (dow & 0xff);
  112.  
  113.   msg[9] = ((((hour >= 12) ? ToBCD(hour) + 0x40 : ToBCD(hour)) & 0xff) << 24) |
  114.            ((ToBCD(minute) & 0xff) << 16) |
  115.            ((ToBCD(second) & 0xff) << 8);
  116.  
  117.   msg[10] = 0;
  118.   msg[11] = 0;
  119.  
  120.   msg[12] = SwapEndianess32(heldButtons ^ ButtonMask);
  121.  
  122.   msg[13] = 0x80000000;
  123.   msg[14] = 0x00000000;
  124.   msg[15] = 0x000001A0; // 416
  125. }
  126.  
  127.  
  128. static char *
  129. LoadProgramSourceFromFile(const char *filename)
  130. {
  131.   struct stat statbuf;
  132.   FILE        *fh;
  133.   char        *source;
  134.  
  135.   fh = fopen(filename, "r");
  136.   if (fh == 0)
  137.     return 0;
  138.  
  139.   stat(filename, &statbuf);
  140.   source = (char *) malloc(statbuf.st_size + 1);
  141.   fread(source, statbuf.st_size, 1, fh);
  142.   source[statbuf.st_size] = '\0';
  143.  
  144.   return source;
  145. }
  146.  
  147.  
  148. cl_device_id      ComputeDeviceId[2];
  149. cl_command_queue  ComputeCommands;
  150. cl_context        ComputeContext;
  151. cl_program        ComputeProgram;
  152. cl_kernel         ComputeKernel;
  153. cl_mem            base_message_buffer;
  154. cl_mem            output_buffer;
  155.  
  156. #define TIMER0_SEARCH_SIZE 0x1000
  157. #define VCOUNT_SEARCH_SIZE 0x100
  158. #define VFRAME_SEARCH_SIZE 0x8
  159.  
  160. int DoOneSearch(uint32_t *msg, uint64_t preSeed, uint32_t vframeBase)
  161. {
  162.   int       err = 0;
  163.  
  164.   uint64_t  t0 = GetCurrentTime();
  165.  
  166.  
  167.   // Fill the input buffer with the host generated message
  168.   //
  169.   err = clEnqueueWriteBuffer(ComputeCommands, base_message_buffer, CL_TRUE, 0,
  170.                              sizeof(uint32_t[16]), msg, 0, NULL, NULL);
  171.   if (err != CL_SUCCESS)
  172.   {
  173.       std::cout << "Error: Failed to write to source array!" << std::endl;
  174.       return EXIT_FAILURE;
  175.   }
  176.  
  177.   // Set the arguments to our compute kernel
  178.   //
  179.   err = 0;
  180.   err  = clSetKernelArg(ComputeKernel, 0, sizeof(cl_mem), &base_message_buffer);
  181.   err |= clSetKernelArg(ComputeKernel, 1, sizeof(uint64_t), &preSeed);
  182.   err |= clSetKernelArg(ComputeKernel, 2, sizeof(cl_mem), &output_buffer);
  183.   if (err != CL_SUCCESS)
  184.   {
  185.     std::cout << "Error: Failed to set kernel arguments! " << err << std::endl;
  186.     exit(1);
  187.   }
  188.  
  189.   // Get the maximum work group size for executing the kernel on the device
  190.   //
  191.   size_t  local[3] = { 0, 1, 1 };
  192.   err = clGetKernelWorkGroupInfo
  193.           (ComputeKernel, ComputeDeviceId[0], CL_KERNEL_WORK_GROUP_SIZE,
  194.            sizeof(local[0]), &local[0], NULL);
  195.   if (err != CL_SUCCESS)
  196.   {
  197.     std::cout << "Error: Failed to retrieve kernel work group info! " << err << std::endl;
  198.     exit(1);
  199.   }
  200.   //std::cout << "Kernel work group size = " << local[0] << std::endl;
  201.   local[0] = GetMaxPowerOf2(local[0]);
  202.   //std::cout << "Work group size = " << local[0] << std::endl;
  203.  
  204.   // Execute the kernel over the entire range of our 1d input data set
  205.   // using the maximum number of work group items for this device
  206.   //
  207.   size_t  global[3] = { TIMER0_SEARCH_SIZE >> 3, VCOUNT_SEARCH_SIZE, VFRAME_SEARCH_SIZE };
  208.   //std::cout << "Global work size = " << global[0] << std::endl;
  209.   err = clEnqueueNDRangeKernel(ComputeCommands, ComputeKernel, 3, NULL, global,
  210.                                local, 0, NULL, NULL);
  211.   if (err != CL_SUCCESS)
  212.   {
  213.     std::cout << "Error: Failed to execute kernel! " << err << std::endl;
  214.     return EXIT_FAILURE;
  215.   }
  216.  
  217.   // Wait for the command commands to get serviced before reading back results
  218.   //
  219.   err = clFinish(ComputeCommands);
  220.   if (err != CL_SUCCESS)
  221.   {
  222.     std::cout << "Error: Failed to finish commands! " << err << std::endl;
  223.     return EXIT_FAILURE;
  224.   }
  225.  
  226.   // Read back the results from the device to verify the output
  227.   //
  228.   uint8_t  *result = new uint8_t[TIMER0_SEARCH_SIZE * VCOUNT_SEARCH_SIZE * VFRAME_SEARCH_SIZE];
  229.   err = clEnqueueReadBuffer(ComputeCommands, output_buffer, CL_TRUE, 0,
  230.                             TIMER0_SEARCH_SIZE * VCOUNT_SEARCH_SIZE * VFRAME_SEARCH_SIZE,
  231.                             result, 0, NULL, NULL);  
  232.   if (err != CL_SUCCESS)
  233.   {
  234.     std::cout << "Error: Failed to read output array! " << err << std::endl;
  235.     return EXIT_FAILURE;
  236.   }
  237.  
  238.   uint8_t  *rptr = result;
  239.   for (int vf = 0; vf < VFRAME_SEARCH_SIZE; ++vf)
  240.   {
  241.     for (int vc = 0; vc < VCOUNT_SEARCH_SIZE; ++vc)
  242.     {
  243.       for (int tmr0 = 0; tmr0 < TIMER0_SEARCH_SIZE; ++tmr0)
  244.       {
  245.         if (*rptr++ != 0)
  246.         {
  247.           std::cout << std::setfill('0') << std::hex;
  248.           std::cout << "\n*****************   Match found at timer0 = " << (tmr0 + 0x1000)
  249.                     << ", vcount = " << vc
  250.                     << ", vframe = " << (vframeBase | vf) << std::endl;
  251.           std::cout << std::dec;
  252.         }
  253.       }
  254.     }
  255.   }
  256.  
  257.   std::cout << std::dec;
  258.  
  259.   uint64_t  t1 = GetCurrentTime();
  260.  
  261.   double t = SubtractTimeInSec(t1, t0);
  262.   std::cout << "Exec Time:  " << (1000.0 * t) << " ms";
  263.  
  264.   delete[] result;
  265.  
  266.   return EXIT_SUCCESS;
  267. }
  268.  
  269.  
  270. int DoFullSearch()
  271. {
  272.   uint32_t  year = 2012;
  273.   uint32_t  month = 10;
  274.   uint32_t  day = 12;
  275.   uint32_t  dayOfWeek = 5;
  276.   uint32_t  hour = 20;
  277.   uint32_t  min = 0;
  278.   uint32_t  second = 23;
  279.   uint32_t  macHigh = 0x8C56C5;
  280.   uint32_t  macLow = 0xE57B30;
  281.   uint64_t  fullseed = 0x69E235BA6219A09CULL;
  282.  
  283.   uint64_t  preSeed = SwapEndianess64((fullseed * 0xDEDCEDAE9638806DULL) + 0x9B1AE6E9A384E6F9ULL);
  284.  
  285.   std::cout << std::setfill('0');
  286.   std::cout << std::hex;
  287.   std::cout << "Searching for pre-seed " << std::setw(16) << preSeed
  288.             << " of actual seed " << std::setw(16) << fullseed << std::endl;
  289.  
  290.   uint32_t  msg[16];
  291.   uint32_t  nazo = 0x027a6a90;
  292.  
  293.   for (uint32_t t = 2; t < 9; ++t)
  294.   {
  295.     while (nazo > 0x027A3000)
  296.     {
  297.       std::cout << std::hex;
  298.       std::cout << "\nNazo = 0x" << nazo << ": Second = ";
  299.       std::cout << std::dec << (second - t);
  300.      
  301.       std::cout << " vframeBase = 0x0, ";
  302.       MakeMessage(msg, nazo, HardResetGxStat, 0, 0, 0, macLow, macHigh,
  303.                   year, month, day, dayOfWeek, hour, min, second - t, 0);
  304.       std::cout.flush();
  305.       if (DoOneSearch(msg, preSeed, 0) != EXIT_SUCCESS)
  306.       {
  307.         std::cerr << "Error encountered - terminating" << std::endl;
  308.         return EXIT_FAILURE;
  309.       }
  310.      
  311.       std::cout << "  -  vframeBase = 0x8, ";
  312.       std::cout.flush();
  313.       MakeMessage(msg, nazo, HardResetGxStat, 0, 8, 0, macLow, macHigh,
  314.                   year, month, day, dayOfWeek, hour, min, second - t, 0);
  315.       if (DoOneSearch(msg, preSeed, 8) != EXIT_SUCCESS)
  316.       {
  317.         std::cerr << "Error encountered - terminating" << std::endl;
  318.         return EXIT_FAILURE;
  319.       }
  320.       std::cout.flush();
  321.      
  322.       nazo -= 0x20;
  323.     }
  324.    
  325.     nazo = 0x027A8010;
  326.   }
  327.  
  328.   return EXIT_SUCCESS;
  329. }
  330.  
  331.  
  332. int main (int argc, char * const argv[])
  333. {
  334.   int       err = 0;
  335.  
  336.   // Connect to a GPU compute device
  337.   //
  338.   err = clGetDeviceIDs(NULL, CL_DEVICE_TYPE_CPU, 1, ComputeDeviceId, NULL);
  339.   if (err != CL_SUCCESS)
  340.   {
  341.       std::cout << "Error: Failed to locate a compute device!" << std::endl;
  342.       return EXIT_FAILURE;
  343.   }
  344.   //ComputeDeviceId[0] = ComputeDeviceId[1];
  345.  
  346.   size_t  returned_size;
  347.   cl_char vendor_name[1024] = {0};
  348.   cl_char device_name[1024] = {0};
  349.   err = clGetDeviceInfo(ComputeDeviceId[0], CL_DEVICE_VENDOR, sizeof(vendor_name),
  350.                         vendor_name, &returned_size);
  351.   err|= clGetDeviceInfo(ComputeDeviceId[0], CL_DEVICE_NAME, sizeof(device_name),
  352.                         device_name, &returned_size);
  353.   if (err != CL_SUCCESS)
  354.   {
  355.       std::cout << "Error: Failed to retrieve device info!" << std::endl;
  356.       return EXIT_FAILURE;
  357.   }
  358.  
  359.   std::cout << SEPARATOR << std::endl;
  360.   std::cout << "Connecting to " << vendor_name << ' ' << device_name << std::endl;
  361.  
  362.   // Load the compute program from disk into a cstring buffer
  363.   //
  364.   std::cout << SEPARATOR << std::endl;
  365.   std::cout << "Loading program '../../sha1.cl' ..." << std::endl;
  366.   std::cout << SEPARATOR << std::endl;
  367.  
  368.   char  *source = LoadProgramSourceFromFile("../../sha1.cl");
  369.  
  370.   // Create a compute ComputeContext
  371.   //
  372.   ComputeContext = clCreateContext(0, 1, &ComputeDeviceId[0], NULL, NULL, &err);
  373.   if (!ComputeContext)
  374.   {
  375.       std::cout << "Error: Failed to create a compute ComputeContext!" << std::endl;
  376.       return EXIT_FAILURE;
  377.   }
  378.  
  379.   // Create a command queue
  380.   //
  381.   ComputeCommands = clCreateCommandQueue(ComputeContext, ComputeDeviceId[0], 0, &err);
  382.   if (!ComputeCommands)
  383.   {
  384.       std::cout << "Error: Failed to create a command ComputeCommands!" << std::endl;
  385.       return EXIT_FAILURE;
  386.   }
  387.  
  388.   // Create the compute program from the source buffer
  389.   //
  390.   ComputeProgram = clCreateProgramWithSource(ComputeContext, 1,
  391.                                              (const char **) &source,
  392.                                              NULL, &err);
  393.   if (!ComputeProgram || err != CL_SUCCESS)
  394.   {
  395.       std::cout << "Error: Failed to create compute program!" << std::endl;
  396.       return EXIT_FAILURE;
  397.   }
  398.  
  399.   // Build the program executable
  400.   //
  401.   err = clBuildProgram(ComputeProgram, 0, NULL, NULL, NULL, NULL);
  402.   if (err != CL_SUCCESS)
  403.   {
  404.       size_t length;
  405.       char build_log[2048];
  406.       std::cout << "Error: Failed to build program executable!" << std::endl;
  407.       clGetProgramBuildInfo(ComputeProgram, ComputeDeviceId[0], CL_PROGRAM_BUILD_LOG, sizeof(build_log), build_log, &length);
  408.       std::cout << build_log << std::endl;
  409.       return EXIT_FAILURE;
  410.   }
  411.  
  412.   // Create the compute kernel from within the program
  413.   //
  414.   ComputeKernel = clCreateKernel(ComputeProgram, "sha1_v0", &err);
  415.   if (!ComputeKernel || err != CL_SUCCESS)
  416.   {
  417.       std::cout << "Error: Failed to create compute kernel!" << std::endl;
  418.       return EXIT_FAILURE;
  419.   }
  420.  
  421.   // Create the input buffer on the device
  422.   //
  423.   size_t buffer_size = sizeof(uint32_t) * 16;
  424.   base_message_buffer = clCreateBuffer(ComputeContext, CL_MEM_READ_ONLY, buffer_size, NULL, NULL);
  425.   if (!base_message_buffer)
  426.   {
  427.       std::cout << "Error: Failed to allocate base message buffer on device!" << std::endl;
  428.       return EXIT_FAILURE;
  429.   }
  430.  
  431.   // Create the output buffer on the device
  432.   //
  433.   buffer_size = TIMER0_SEARCH_SIZE * VCOUNT_SEARCH_SIZE * VFRAME_SEARCH_SIZE;
  434.   output_buffer = clCreateBuffer(ComputeContext, CL_MEM_WRITE_ONLY, buffer_size, NULL, &err);
  435.   if (!output_buffer)
  436.   {
  437.       std::cout << "Error: Failed to allocate result buffer on device! " << err << std::endl;
  438.       return EXIT_FAILURE;
  439.   }
  440.  
  441.   DoFullSearch();
  442.  
  443.   // Shutdown and cleanup
  444.   //
  445.   clReleaseKernel(ComputeKernel);
  446.   clReleaseProgram(ComputeProgram);
  447.   clReleaseMemObject(base_message_buffer);
  448.   clReleaseMemObject(output_buffer);
  449.   clReleaseCommandQueue(ComputeCommands);
  450.   clReleaseContext(ComputeContext);
  451.      
  452.   return 0;
  453. }
RAW Paste Data
We use cookies for various purposes including analytics. By continuing to use Pastebin, you agree to our use of cookies as described in the Cookies Policy. OK, I Understand
 
Top