Pastebin launched a little side project called VERYVIRAL.com, check it out ;-) Want more features on Pastebin? Sign Up, it's FREE!
Guest

Untitled

By: a guest on Oct 29th, 2012  |  syntax: C  |  size: 1.94 KB  |  views: 53  |  expires: Never
download  |  raw  |  embed  |  report abuse  |  print
Text below is selected. Please press Ctrl+C to copy to your clipboard. (⌘+C on Mac)
  1. // The kernel (this is the entrypoint of GPU code)
  2. // Loads the 64-byte word to be hashed from g   lobal to shared memory and calls the calculation routine
  3. __global__ void md5_calc(char *gwords, char *paddedWords3, uint *hash, int realthreads, int msg_size, int size_hash, uint launch){
  4.         int linidx;
  5.         uint a;
  6.         uint b = 0;
  7.         uint c;
  8.         uint d;
  9.  
  10.         char msg[64] = {0};
  11.  
  12.         // assuming blockDim.y = 1 and threadIdx.y = 0, always
  13.         int iblock= blockIdx.x + blockIdx.y * gridDim.x;
  14.         linidx = threadIdx.x + iblock * blockDim.x;
  15.  
  16.         // this check slows down the code by ~0.4% (measured)
  17.         if(linidx >= realthreads){
  18.                 return;
  19.         }
  20.  
  21.         // load the dictionary word for this thread
  22.         //char *word = &memory[0] + threadIdx.x * msg_size;
  23.  
  24.         for(a = 0; a < msg_size; a++){
  25.                 //word[(a + threadIdx.x) % msg_size] = gwords[linidx * msg_size + a];
  26.                 //msg[a] = word[(a + threadIdx.x) % msg_size];
  27.                 msg[a] = gwords[linidx * msg_size + a];
  28.  
  29.                 //end of the password
  30.                 if(msg[a] == 0){
  31.                         msg[a] = 0x80;
  32.                         b = a;
  33.  
  34.                         break;
  35.                 }
  36.         }
  37.  
  38.         if(!b){
  39.                 msg[msg_size] = 0x80;
  40.        
  41.                 msg[56] = msg_size * 8;
  42.         }
  43.         else{
  44.                 msg[56] = b * 8;
  45.         }
  46.  
  47.  
  48.         for(a = 0; a < 64; a++){
  49.                 paddedWords3[linidx * 64 + a] = msg[a];
  50.         }
  51.  
  52.         paddedWords3[linidx * 64 + 63] = linidx;
  53.         paddedWords3[linidx * 64 + 62] = launch;
  54.         paddedWords3[linidx * 64 + 61] = (launch * realthreads) + linidx + 1;
  55.  
  56.         // compute MD5 hash
  57.         md5_gpu((uint *)msg, a, b, c, d);
  58.         //md5_gpu((uint *)msg1, a1, b1, c1, d1);
  59.  
  60.         // return the hash converted to number
  61.         // all this work around is because of different edian order
  62.         if(size_hash == 1){
  63.                 hash[(linidx)] = ((a & 240) >> 4);
  64.         }
  65.         else{
  66.                 if(size_hash == 2){
  67.                         hash[(linidx)] = a & 255;
  68.                 }
  69.                 else{
  70.                         if(size_hash == 3){
  71.                                 hash[(linidx)] = ((a & 61440) >> 12) + ((a & 240) << 4) + ((a & 15) << 4);
  72.                         }
  73.                         else{
  74.                                 if(size_hash == 4){
  75.                                         hash[(linidx)] = ((a & 61440) >> 8) + ((a & 3840) >> 8) + ((a & 240) << 8) + ((a & 15) << 8);
  76.                                 }
  77.                         }
  78.                 }
  79.         }
  80. }