/*

 this program compute md5 on cuda device and try to discover the clear
 text by brute force.

 Copyright (C) 2008 Luigi Tarenga <luigi dot tarenga at gmail dot com>

 This program is free software; you can redistribute it and/or
 modify it under the terms of the GNU General Public License
 as published by the Free Software Foundation; either version 2
 of the License, or (at your option) any later version.

 This program is distributed in the hope that it will be useful,
 but WITHOUT ANY WARRANTY; without even the implied warranty of
 MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
 GNU General Public License for more details.

 You should have received a copy of the GNU General Public License
 along with this program; if not, write to the Free Software
 Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301, USA.

 This code take inspiration from wikipedia MD5 explanation and code from L. Peter Deutsch 

*/

#include<stdlib.h>
#include<stdio.h>
#include<string.h>
#include <arpa/inet.h>

#define BLOCK_SIZE 96
#define GRID_SIZE 96

#define ROTATE_LEFT(x, n) (((x) << (n)) | ((x) >> (32 - (n))))
#define F(x, y ,z) ((z) ^ ((x) & ((y) ^ (z)))) /*(((x) & (y)) | (~(x) & (z)))*/
#define G(x, y, z) ((y) ^ ((z) & ((x) ^ (y)))) /*(((x) & (z)) | ((y) & ~(z)))*/
#define H(x, y, z) ((x) ^ (y) ^ (z))
#define I(x, y, z) ((y) ^ ((x) | ~(z)))
#define R1OP(a, b, c, d, idx, s, K) t = a + F(b,c,d) + intmess[idx] + K; a = ROTATE_LEFT(t, s) + b
#define R2OP(a, b, c, d, idx, s, K) t = a + G(b,c,d) + intmess[idx] + K; a = ROTATE_LEFT(t, s) + b
#define R3OP(a, b, c, d, idx, s, K) t = a + H(b,c,d) + intmess[idx] + K; a = ROTATE_LEFT(t, s) + b
#define R4OP(a, b, c, d, idx, s, K) t = a + I(b,c,d) + intmess[idx] + K; a = ROTATE_LEFT(t, s) + b


/* formula per generare le costanti
for i from 0 to 63
    K := floor(abs(sin(i + 1)) × (2 pow 32))
*/

__global__ void cudaMD5(unsigned char *mess, unsigned int len, uint4 digest){
  unsigned int intmess[16];
  unsigned char *charmess;
  unsigned int i,j;
  unsigned int a;
  unsigned int b;
  unsigned int c;
  unsigned int d;
  unsigned int t;


  charmess=(unsigned char *)intmess;

  /* trick to save 4 sum on the gpu */
  digest.x-=0x67452301;
  digest.y-=0xefcdab89;
  digest.z-=0x98badcfe;
  digest.w-=0x10325476;

  for(i=0;i<16;i++){ intmess[i]=0; }

  charmess[0]=(unsigned char)(blockIdx.x+32);
  charmess[1]=(unsigned char)(threadIdx.x+32);
  for(i=2;i<len;i++){
    charmess[i]=32;
  }
  /* padding: aggiungo 1 bit a 1 al messaggio e 0 fino a 512-64 bit. l'ultimo long int (64bit) contiene il numero di bit del messaggio unpadded */
  charmess[len]=128;

  intmess[14]=len*8;
  intmess[15]=0;

  while(charmess[len-1]<127){

    /* inizializzazione  */
    a=0x67452301;
    b=0xefcdab89; /* (~0) ^ 0x10325476 */
    c=0x98badcfe; /* (~0) ^ 0x67452301 */
    d=0x10325476;

    R1OP(a, b, c, d,  0,  7, 0xd76aa478);
    R1OP(d, a, b, c,  1, 12, 0xe8c7b756);
    R1OP(c, d, a, b,  2, 17, 0x242070db);
    R1OP(b, c, d, a,  3, 22, 0xc1bdceee);
    R1OP(a, b, c, d,  4,  7, 0xf57c0faf);
    R1OP(d, a, b, c,  5, 12, 0x4787c62a);
    R1OP(c, d, a, b,  6, 17, 0xa8304613);
    R1OP(b, c, d, a,  7, 22, 0xfd469501);
    R1OP(a, b, c, d,  8,  7, 0x698098d8);
    R1OP(d, a, b, c,  9, 12, 0x8b44f7af);
    R1OP(c, d, a, b, 10, 17, 0xffff5bb1);
    R1OP(b, c, d, a, 11, 22, 0x895cd7be);
    R1OP(a, b, c, d, 12,  7, 0x6b901122);
    R1OP(d, a, b, c, 13, 12, 0xfd987193);
    R1OP(c, d, a, b, 14, 17, 0xa679438e);
    R1OP(b, c, d, a, 15, 22, 0x49b40821);

    R2OP(a, b, c, d,  1,  5, 0xf61e2562);
    R2OP(d, a, b, c,  6,  9, 0xc040b340);
    R2OP(c, d, a, b, 11, 14, 0x265e5a51);
    R2OP(b, c, d, a,  0, 20, 0xe9b6c7aa);
    R2OP(a, b, c, d,  5,  5, 0xd62f105d);
    R2OP(d, a, b, c, 10,  9, 0x02441453);
    R2OP(c, d, a, b, 15, 14, 0xd8a1e681);
    R2OP(b, c, d, a,  4, 20, 0xe7d3fbc8);
    R2OP(a, b, c, d,  9,  5, 0x21e1cde6);
    R2OP(d, a, b, c, 14,  9, 0xc33707d6);
    R2OP(c, d, a, b,  3, 14, 0xf4d50d87);
    R2OP(b, c, d, a,  8, 20, 0x455a14ed);
    R2OP(a, b, c, d, 13,  5, 0xa9e3e905);
    R2OP(d, a, b, c,  2,  9, 0xfcefa3f8);
    R2OP(c, d, a, b,  7, 14, 0x676f02d9);
    R2OP(b, c, d, a, 12, 20, 0x8d2a4c8a);

    R3OP(a, b, c, d,  5,  4, 0xfffa3942);
    R3OP(d, a, b, c,  8, 11, 0x8771f681);
    R3OP(c, d, a, b, 11, 16, 0x6d9d6122);
    R3OP(b, c, d, a, 14, 23, 0xfde5380c);
    R3OP(a, b, c, d,  1,  4, 0xa4beea44);
    R3OP(d, a, b, c,  4, 11, 0x4bdecfa9);
    R3OP(c, d, a, b,  7, 16, 0xf6bb4b60);
    R3OP(b, c, d, a, 10, 23, 0xbebfbc70);
    R3OP(a, b, c, d, 13,  4, 0x289b7ec6);
    R3OP(d, a, b, c,  0, 11, 0xeaa127fa);
    R3OP(c, d, a, b,  3, 16, 0xd4ef3085);
    R3OP(b, c, d, a,  6, 23, 0x04881d05);
    R3OP(a, b, c, d,  9,  4, 0xd9d4d039);
    R3OP(d, a, b, c, 12, 11, 0xe6db99e5);
    R3OP(c, d, a, b, 15, 16, 0x1fa27cf8);
    R3OP(b, c, d, a,  2, 23, 0xc4ac5665);

    R4OP(a, b, c, d,  0,  6, 0xf4292244);
    R4OP(d, a, b, c,  7, 10, 0x432aff97);
    R4OP(c, d, a, b, 14, 15, 0xab9423a7);
    R4OP(b, c, d, a,  5, 21, 0xfc93a039);
    R4OP(a, b, c, d, 12,  6, 0x655b59c3);
    R4OP(d, a, b, c,  3, 10, 0x8f0ccc92);
    R4OP(c, d, a, b, 10, 15, 0xffeff47d);
    R4OP(b, c, d, a,  1, 21, 0x85845dd1);
    R4OP(a, b, c, d,  8,  6, 0x6fa87e4f);
    R4OP(d, a, b, c, 15, 10, 0xfe2ce6e0);
    R4OP(c, d, a, b,  6, 15, 0xa3014314);
    R4OP(b, c, d, a, 13, 21, 0x4e0811a1);
    R4OP(a, b, c, d,  4,  6, 0xf7537e82);
    R4OP(d, a, b, c, 11, 10, 0xbd3af235);
    R4OP(c, d, a, b,  2, 15, 0x2ad7d2bb);
    R4OP(b, c, d, a,  9, 21, 0xeb86d391);

    if(a==digest.x && b==digest.y && c==digest.z && d==digest.w){

#pragma unroll 1
      for(j=0;j<64;j++){
        mess[j]=charmess[j];
      }
      mess[len]=0;

    }

    /* generate a new string to use */
#pragma unroll 1
    for(i=2;i<len;i++){
      if(charmess[i]<127){
        charmess[i]++;
        i=len;
      } else {
        charmess[i]=32;
      }
    }

  } /*while */
}


int main(int argc, char **argv){
  unsigned char message[65];
  unsigned char *dev_ptr;
  uint4 d;
  cudaError_t err;
  unsigned int i;
  dim3 DimGrid(GRID_SIZE);
  dim3 DimBlock(BLOCK_SIZE);

  for(i=0;i<65;i++){
    message[i]=0;
  }

  if(argc<=1){
    printf("\nyou must specify a digest to crack\n");
    exit(1);
  }
  if(strlen(&argv[1][0])<32){
    printf("invalid digest\n");
    exit(1);
  }

  /* leggo il digest da crackare e lo salvo in 4 int convertendo da big endian a little endian */
  sscanf(argv[1],"%08x%08x%08x%08x",&d.x,&d.y,&d.z,&d.w); 
  d.x=ntohl(d.x); d.y=ntohl(d.y); d.z=ntohl(d.z); d.w=ntohl(d.w);
  printf("\ncracking %08x %08x %08x %08x\n",d.x,d.y,d.z,d.w);

  err=cudaMalloc((void **)&dev_ptr,64); 
  if( cudaSuccess != err ){
    printf("errore in cudaMalloc\n");
    exit(1);
  }

  cudaMD5<<<DimGrid , DimBlock>>>(dev_ptr,4,d);
  err=cudaThreadSynchronize();
  if( cudaSuccess != err ){
    printf("errore in cudaThreadSynchronize\n");
    exit(1);
  }
 
  err=cudaMemcpy((void *)message,(void *)dev_ptr,64,cudaMemcpyDeviceToHost);
  if( cudaSuccess != err ){
    printf("errore in cudaMemcpy\n");
    exit(1);
  }
  err=cudaFree((void *)dev_ptr);
  if( cudaSuccess != err ){
    printf("errore in cudaFree\n");
    exit(1);
  }

  printf("la stringa in chiaro è: \"%s\"\n",message);
  for(i=0;i<64;i++){
    printf("%02x ",message[i]);
  }
  printf("\n");
  exit(0);
}