#include<math.h>
#include<omp.h>
#include<time.h>
#include<iostream>
#include<CL/cl.h>

#define MAX_SALT_LENGTH 44

#define INIT_MD4_A                  0x67452301
#define INIT_MD4_B                  0xefcdab89
#define INIT_MD4_C                  0x98badcfe
#define INIT_MD4_D                  0x10325476

#define SQRT_2                      0x5a827999
#define SQRT_3                      0x6ed9eba1

unsigned int NUM_KEYS;


void md4_crypt(unsigned int *buffer, unsigned int *hash)
{
    unsigned int a;
    unsigned int b;
    unsigned int c;
    unsigned int d;
 
    // round 1
    a = 0xFFFFFFFF  +  buffer[0]; a = (a << 3 ) | (a >> 29);
    d = INIT_MD4_D + (INIT_MD4_C ^ (a & 0x77777777)) + buffer[1]; d = (d << 7 ) | (d >> 25);
    c = INIT_MD4_C + (INIT_MD4_B ^ (d & (a ^ INIT_MD4_B))) + buffer[2]; c = (c << 11) | (c >> 21);
    b = INIT_MD4_B + (a ^ (c & (d ^ a))) + buffer[3]; b = (b << 19) | (b >> 13); 
 
    a += (d ^ (b & (c ^ d)))  +  buffer[4];  a = (a << 3 ) | (a >> 29);
    d += (c ^ (a & (b ^ c)))  +  buffer[5];  d = (d << 7 ) | (d >> 25);
    c += (b ^ (d & (a ^ b)))  +  buffer[6];  c = (c << 11) | (c >> 21);
    b += (a ^ (c & (d ^ a)))  +  buffer[7];  b = (b << 19) | (b >> 13);
 
    a += (d ^ (b & (c ^ d)))  + buffer[8] ;  a = (a << 3 ) | (a >> 29);
    d += (c ^ (a & (b ^ c)))  + buffer[9] ;  d = (d << 7 ) | (d >> 25);
    c += (b ^ (d & (a ^ b)))  + buffer[10];  c = (c << 11) | (c >> 21);
    b += (a ^ (c & (d ^ a)))  + buffer[11];  b = (b << 19) | (b >> 13);
 
    a += (d ^ (b & (c ^ d)))  + buffer[12]; a = (a << 3 ) | (a >> 29);
    d += (c ^ (a & (b ^ c)))  + buffer[13]; d = (d << 7 ) | (d >> 25);
    c += (b ^ (d & (a ^ b)))  + buffer[14]; c = (c << 11) | (c >> 21);
    b += (a ^ (c & (d ^ a)))  + buffer[15]; b = (b << 19) | (b >> 13);
 
    // round 2
    a += ((b & (c | d)) | (c & d)) + buffer[0]  + SQRT_2; a = (a<<3 ) | (a>>29);
    d += ((a & (b | c)) | (b & c)) + buffer[4]  + SQRT_2; d = (d<<5 ) | (d>>27);
    c += ((d & (a | b)) | (a & b)) + buffer[8]  + SQRT_2; c = (c<<9 ) | (c>>23);
    b += ((c & (d | a)) | (d & a)) + buffer[12] + SQRT_2; b = (b<<13) | (b>>19);
 
    a += ((b & (c | d)) | (c & d)) + buffer[1]  + SQRT_2; a = (a<<3 ) | (a>>29);
    d += ((a & (b | c)) | (b & c)) + buffer[5]  + SQRT_2; d = (d<<5 ) | (d>>27);
    c += ((d & (a | b)) | (a & b)) + buffer[9]  + SQRT_2; c = (c<<9 ) | (c>>23);
    b += ((c & (d | a)) | (d & a)) + buffer[13] + SQRT_2; b = (b<<13) | (b>>19);
 
    a += ((b & (c | d)) | (c & d)) + buffer[2]  + SQRT_2; a = (a<<3 ) | (a>>29);
    d += ((a & (b | c)) | (b & c)) + buffer[6]  + SQRT_2; d = (d<<5 ) | (d>>27);
    c += ((d & (a | b)) | (a & b)) + buffer[10] + SQRT_2; c = (c<<9 ) | (c>>23);
    b += ((c & (d | a)) | (d & a)) + buffer[14] + SQRT_2; b = (b<<13) | (b>>19);
 
    a += ((b & (c | d)) | (c & d)) + buffer[3]  + SQRT_2; a = (a<<3 ) | (a>>29);
    d += ((a & (b | c)) | (b & c)) + buffer[7]  + SQRT_2; d = (d<<5 ) | (d>>27);
    c += ((d & (a | b)) | (a & b)) + buffer[11] + SQRT_2; c = (c<<9 ) | (c>>23);
    b += ((c & (d | a)) | (d & a)) + buffer[15] + SQRT_2; b = (b<<13) | (b>>19);         
 
    // round 3
    a += (d ^ c ^ b) + buffer[0]  +  SQRT_3; a = (a << 3 ) | (a >> 29);
    d += (c ^ b ^ a) + buffer[8]  +  SQRT_3; d = (d << 9 ) | (d >> 23);
    c += (b ^ a ^ d) + buffer[4]  +  SQRT_3; c = (c << 11) | (c >> 21);
    b += (a ^ d ^ c) + buffer[12] +  SQRT_3; b = (b << 15) | (b >> 17);
 
    a += (d ^ c ^ b) + buffer[2]  +  SQRT_3; a = (a << 3 ) | (a >> 29);
    d += (c ^ b ^ a) + buffer[10] +  SQRT_3; d = (d << 9 ) | (d >> 23);
    c += (b ^ a ^ d) + buffer[6]  +  SQRT_3; c = (c << 11) | (c >> 21);
    b += (a ^ d ^ c) + buffer[14] +  SQRT_3; b = (b << 15) | (b >> 17);
 
    a += (d ^ c ^ b) + buffer[1]  +  SQRT_3; a = (a << 3 ) | (a >> 29);
    d += (c ^ b ^ a) + buffer[9]  +  SQRT_3; d = (d << 9 ) | (d >> 23);
    c += (b ^ a ^ d) + buffer[5]  +  SQRT_3; c = (c << 11) | (c >> 21);
    b += (a ^ d ^ c) + buffer[13] +  SQRT_3; b = (b << 15) | (b >> 17);
 
    a += (d ^ c ^ b) + buffer[3]  +  SQRT_3; a = (a << 3 ) | (a >> 29);
 
    d += (c ^ b ^ a) + buffer[11] +  SQRT_3; d = (d << 9 ) | (d >> 23);
    c += (b ^ a ^ d) + buffer[7]  +  SQRT_3; c = (c << 11) | (c >> 21);
    b += (a ^ d ^ c) + buffer[15] +  SQRT_3; b = (b << 15) | (b >> 17);
 
    hash[0] = a + INIT_MD4_A;
    hash[1] = b + INIT_MD4_B;
    hash[2] = c + INIT_MD4_C;
    hash[3] = d + INIT_MD4_D;
}

unsigned char *byte2hexstring(unsigned char * byte, unsigned int len) {
    unsigned int i;
    unsigned char *hexstring;
 
    hexstring =(unsigned char*) malloc(len * 2 + 1);
    memset(hexstring,0, 2 * len + 1);
 
    for (i = 0; i < len; i++)
        sprintf((char*)&hexstring[2 * i], "%02x", byte[i]);
 
    return hexstring;
}

//Read a program containing kernels
char* read_kernel(const char *file_name,bool display_source,long *filelength)
{ FILE *fp;
if(!(fp=fopen(file_name,"r+"))) {std::cout<<"Cannot open file.Wrong file name or file doesn't exsist.";return NULL;}
fseek(fp,2,SEEK_END);
fputs(" ",fp);
fseek(fp,0,SEEK_SET);
  long i=0;
  *filelength=0;
  while(!feof(fp))
     {   getc(fp);(*filelength)++;   }
  fclose(fp);
  if(!(fp=fopen(file_name,"r+"))) {std::cout<<"Cannot open file.Wrong file name or file doesn't exsist.";return NULL;}
  char  *kernel_source=new char[sizeof(char)*(*filelength+1)] ;
  while(!feof(fp))
     {    kernel_source[i]=getc(fp);
          ++i;
     }
  *filelength=*filelength-1;
  kernel_source[*filelength]='\0';
  if(display_source)
  std::cout<<kernel_source;
  fclose(fp);
  return kernel_source;
}

//Callback function for clBuildProgram();
void __stdcall build_call_back(cl_program prg,void *devid)
{   char *ch;
    cl_device_id did;
	size_t length ;
	memcpy((cl_device_id*) &did,(const void*) devid,sizeof(cl_device_id));
    clGetProgramBuildInfo(prg,did,CL_PROGRAM_BUILD_LOG,200,NULL,&length);
	ch=new char[length];
	clGetProgramBuildInfo(prg,did,CL_PROGRAM_BUILD_LOG,length,ch,&length);
	if(length!=1)
	std::cout<<"error string:"<<ch<<std::endl;

}

//Intialize host environment
bool CL_init(cl_platform_id *pltfrmid,cl_device_id *devid,cl_context *cntxt,cl_command_queue *cmdq,const char* file_name,cl_program *prg)
{ cl_int err;
  if(!clGetPlatformIDs(1,pltfrmid,NULL))std::cout<<"PlatformID OK"<<std::endl; else {std::cout<<"PlatformID FAILED"<<std::endl; return false;}
  if(!clGetDeviceIDs(*pltfrmid,CL_DEVICE_TYPE_GPU,1,devid,NULL))std::cout<<"DeviceID OK"<<std::endl; else {std::cout<<"DeviceID FAILED"<<std::endl; return false;}
  cl_context_properties cnxtprop[3]={CL_CONTEXT_PLATFORM,(cl_context_properties) *pltfrmid,0};
  *cntxt=clCreateContext(cnxtprop,1,devid,NULL,NULL,&err);
  if(!err) std::cout<<"Create Context OK"<<std::endl;else {std::cout<<"Create Context FAILED"<<std::endl; return false;}
  *cmdq=clCreateCommandQueue(*cntxt,devid[0],0,&err);
  if(!err) std::cout<<"Create Command Queue OK"<<std::endl; else {std::cout<<"Create Command Queue FAILED"<<std::endl; return false;}
  char *krnl_src;
  long length;
  krnl_src=read_kernel(file_name,0,&length);
  *prg=clCreateProgramWithSource(*cntxt,1,(const char **)&krnl_src,NULL,&err);
  if(!err) std::cout<<"Create Program OK"<<std::endl; else {std::cout<<"Create Program FAILED"<<std::endl; return false;}
  if(!clBuildProgram(*prg,1,&devid[0],NULL,build_call_back,&devid[0]))std::cout<<"Build Program OK"<<std::endl; else {std::cout<<"Build Program FAILED"<<std::endl; return false;}
  return true;
}

void PBKDF2_api(cl_uint *pass_api,cl_uint *salt_api,cl_uint saltlen_api,cl_uint *hash_out_api)
{ 
  cl_platform_id pltfrmid;cl_device_id devid[1];cl_context cntxt;cl_command_queue cmdq;cl_program prg;cl_kernel krnl0; cl_mem pass,salt,hash_out;cl_int err;
  if(!CL_init(&pltfrmid,devid,&cntxt,&cmdq,"PBKDF2.cl",&prg)) return;
  
  pass=clCreateBuffer(cntxt,CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR,4*NUM_KEYS*sizeof(cl_uint),pass_api,&err);
  if((pass==(cl_mem)0))  {std::cout<<"Create Buffer FAILED:"<<std::endl; return;}
  
  salt=clCreateBuffer(cntxt,CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR,MAX_SALT_LENGTH*sizeof(cl_uint)/4,salt_api,&err);
  if((salt==(cl_mem)0)) {std::cout<<"Create Buffer FAILED"<<std::endl; return;}
  
  hash_out=clCreateBuffer(cntxt,CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR,4*NUM_KEYS*sizeof(cl_uint),hash_out_api,&err);
  if((hash_out==(cl_mem)0)) {std::cout<<"Create Buffer FAILED"<<std::endl; return;}
    
  krnl0=clCreateKernel(prg,"PBKDF2",&err) ;
  if(!err) std::cout<<"Create Kernel PBKDF2 OK"<<std::endl; else {std::cout<<"Create Kernel PBKDF2 FAILED"<<std::endl; return ;}
  
  if(clSetKernelArg(krnl0,0,sizeof(cl_mem),&pass)) {std::cout<<"Set Kernel FAILED.krnl0 arg0"<<std::endl; return;}
  if(clSetKernelArg(krnl0,1,sizeof(cl_mem),&salt)) {std::cout<<"Set Kernel FAILED.krnl0 arg1"<<std::endl; return;}
  if(clSetKernelArg(krnl0,2,sizeof(cl_uint),&saltlen_api)) {std::cout<<"Set Kernel FAILED.krnl0 arg2"<<std::endl; return;}
  if(clSetKernelArg(krnl0,3,sizeof(cl_mem),&hash_out))     {std::cout<<"Set Kernel FAILED.krnl0 arg3"<<std::endl; return;}
 
  cl_event evnt;
  size_t N=NUM_KEYS;
  if(clEnqueueNDRangeKernel(cmdq,krnl0,1,NULL,&N,NULL,0,NULL,&evnt)) {std::cout<<"Enqueue Kernel FAILED.krnl0"<<std::endl; return;}
  if(CL_SUCCESS!=clWaitForEvents(1,&evnt)) std::cout<<"SYNC FAILED"<<std::endl;

  if(clEnqueueReadBuffer(cmdq,hash_out,CL_TRUE,0,4*NUM_KEYS*sizeof(cl_uint),hash_out_api, 0, NULL, NULL)) {std::cout<<"Write Read FAILED"<<std::endl; return;} 
     
}

void DCC(unsigned char *salt,unsigned char *username,unsigned int username_len,unsigned char **password,unsigned int *dcc_hash,unsigned int id)
{
    unsigned int i;
    unsigned int buffer[16];
    unsigned int nt_hash[16];
    unsigned int password_len = strlen((char*)password[id]);
    
    memset(nt_hash, 0, 64);
    memset(buffer, 0, 64);
   
    // convert ASCII password to Unicode
     for(i = 0; i < password_len  >> 1; i++)  
        buffer[i] = password[id][2 * i] | (password[id][2 * i + 1] << 16);
 
    // MD4 padding
    if(password_len % 2 == 1)
        buffer[i] = password[id][password_len - 1] | 0x800000;
    else
        buffer[i]=0x80;
 
    // put password length at end of buffer
    buffer[14] = password_len << 4;
 
    // generate MD4 hash of the password (NT hash)
    md4_crypt(buffer, nt_hash);
 
    // concatenate NT hash and the username (salt)
    memcpy((unsigned char *)nt_hash + 16, salt, username_len << 1); 
 
    i = username_len + 8;
 
    // MD4 padding
    if(username_len % 2 == 1)
        nt_hash[i >> 1] = username[username_len - 1] | 0x800000;
    else
        nt_hash[i >> 1] = 0x80;
 
    // put length at end of buffer
    nt_hash[14] = i << 4; 
 
    md4_crypt(nt_hash, (dcc_hash+4*id));
}

int main(int argc, char *argv[])
{  NUM_KEYS=12;
   unsigned int *dcc_hash=(unsigned int *)malloc(4*NUM_KEYS*sizeof(unsigned int));
   unsigned char salt[MAX_SALT_LENGTH];
   unsigned char usrname[]="123456789123";  //Username
   unsigned char **password=(unsigned char**)malloc(NUM_KEYS);
   cl_uint salt_api[MAX_SALT_LENGTH/4];
   cl_uint *dcc_hash_api=(cl_uint*)malloc(16*NUM_KEYS);
   cl_uint *dcc_hash2_api=(cl_uint*)malloc(4*NUM_KEYS*sizeof(cl_uint));
   unsigned int i;
   
   for(i=0;i<NUM_KEYS;i++)
   password[i]=(unsigned char*)malloc(20);
      
   //Generate NUM_KEYS no. of password 
   for(i=0;i<NUM_KEYS;i=i+3)
   {
   password[i]=(unsigned char*)"password12345";
   password[i+1]=(unsigned char*)"qwerty";
   password[i+2]=(unsigned char*)"abcd1234";
   }

   memset(salt,0,MAX_SALT_LENGTH);
  
   unsigned int username_len =strlen((char*)usrname);
   
    // convert ASCII username to Unicode (WideChar)
   if(username_len%2==1)
    for(i = 0; i < (username_len >> 1) + 1; i++)
        ((unsigned int *)salt)[i] = usrname[2 * i] | (usrname[2 * i + 1] << 16);
   else
	   for(i = 0; i < (username_len >> 1) ; i++)
        ((unsigned int *)salt)[i] = usrname[2 * i] | (usrname[2 * i + 1] << 16);
	
	
	for(i=0;i<NUM_KEYS;i++)
	DCC(salt,usrname,username_len,password,dcc_hash,i);
	
	
	memcpy(salt_api,salt,MAX_SALT_LENGTH);
	memcpy(dcc_hash_api,dcc_hash,16*NUM_KEYS);
	
	PBKDF2_api(dcc_hash_api,salt_api,username_len,dcc_hash2_api);
	
	printf("\n  Username           : %s\n", usrname);
	printf("  Password                  MSCash2 hash:\n\n");
	
	for(i=0;i<NUM_KEYS;i++)
	{
	printf("%16s          %s\n",password[i],byte2hexstring((unsigned char*)(dcc_hash2_api+4*i), 16));
	}
	
	
}