From ef871d5603d428351453fd440bbaf81d60ffa274 Mon Sep 17 00:00:00 2001 From: elenago Date: Mon, 20 Nov 2017 20:18:47 +0100 Subject: [PATCH] Option -m (MAC verification) in OpenCL implementation --- OpenCL_version/bitcracker.h | 17 +- OpenCL_version/kernel_attack.cl | 1162 ++++++++++++++++++++++++++++++- OpenCL_version/main.c | 187 ++--- OpenCL_version/opencl_attack.c | 185 ++++- OpenCL_version/utils.c | 22 +- 5 files changed, 1432 insertions(+), 141 deletions(-) diff --git a/OpenCL_version/bitcracker.h b/OpenCL_version/bitcracker.h index ba49ea9..771d469 100755 --- a/OpenCL_version/bitcracker.h +++ b/OpenCL_version/bitcracker.h @@ -58,8 +58,12 @@ #define MAC_SIZE 16 #define NONCE_SIZE 12 #define IV_SIZE 16 + #define VMK_SIZE 60 -#define VMK_DECRYPT_SIZE 16 +#define VMK_HEADER_SIZE 12 +#define VMK_BODY_SIZE 32 +#define VMK_FULL_SIZE 44 + #define DICT_BUFSIZE (50*1024*1024) #define MAX_PLEN 32 @@ -187,6 +191,7 @@ extern int psw_x_thread; extern int tot_psw; extern size_t size_psw; extern int strict_check; +extern int mac_comparison; extern int MAX_PASSWD_SINGLE_KERNEL; extern int DEV_NVIDIA; @@ -207,12 +212,12 @@ extern cl_device_id* cdDevices; // OpenCL device(s) unsigned int * w_block_precomputed(unsigned char * salt); int readFilePassword(char ** buf, int maxNumPsw, FILE *fp); -int parse_data(char *input_hash, - unsigned char ** salt, - unsigned char ** nonce, - unsigned char ** vmk); +int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, unsigned char ** vmk, unsigned char ** mac); +char * opencl_attack(char *dname, unsigned int * w_blocks, + unsigned char * encryptedVMK, + unsigned char * nonce, unsigned char * encryptedMAC, + int gridBlocks); -char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryptedVMK, unsigned char * nonce, int gridBlocks); void setBufferPasswordSize(size_t avail, size_t * passwordBufferSize, int * numPassword); void * Calloc(size_t len, size_t size); diff --git a/OpenCL_version/kernel_attack.cl b/OpenCL_version/kernel_attack.cl index a99340a..d04f915 100755 --- a/OpenCL_version/kernel_attack.cl +++ b/OpenCL_version/kernel_attack.cl @@ -158,6 +158,14 @@ __constant unsigned int TS3[256] = { 0x4141C382U, 0x9999B029U, 0x2D2D775AU, 0x0F0F111EU, 0xB0B0CB7BU, 0x5454FCA8U, 0xBBBBD66DU, 0x16163A2CU }; +unsigned int LOP3LUT_XOR(unsigned int a, unsigned int b, unsigned int c); +unsigned int LOP3LUT_XORAND(unsigned int a, unsigned int b, unsigned int c); +unsigned int LOP3LUT_ANDOR(unsigned int a, unsigned int b, unsigned int c); +void encrypt( + unsigned int k0, unsigned int k1, unsigned int k2, unsigned int k3, unsigned int k4, unsigned int k5, unsigned int k6, unsigned int k7, + unsigned int m0, unsigned int m1, unsigned int m2, unsigned int m3, + unsigned int * output0, unsigned int * output1, unsigned int * output2, unsigned int * output3); + #define MIN(a,b) (((a)<(b))?(a):(b)) #define AES_CTX_LENGTH 256 #define FALSE 0 @@ -166,7 +174,10 @@ __constant unsigned int TS3[256] = { #define MAC_SIZE 16 #define NONCE_SIZE 12 #define IV_SIZE 16 -#define VMK_DECRYPT_SIZE 16 +#define VMK_SIZE 60 +#define VMK_HEADER_SIZE 12 +#define VMK_BODY_SIZE 32 +#define VMK_FULL_SIZE 44 #define DICT_BUFSIZE (50*1024*1024) #define HASH_SIZE 8 //32 @@ -498,7 +509,7 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * first_hash6 = 0x1F83D9AB; first_hash7 = 0x5BE0CD19; - //----------------------------------------------------- FIRST HASH ------------------------------------------------ +//----------------------------------------------------- FIRST HASH ------------------------------------------------ a = 0x6A09E667; b = 0xBB67AE85; c = 0x3C6EF372; @@ -1175,11 +1186,38 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * indexW += SINGLE_BLOCK_W_SIZE; } - schedule0 = IV0 ^ hash0; - schedule1 = IV4 ^ hash1; - schedule2 = IV8 ^ hash2; - schedule3 = IV12 ^ hash3; - + schedule0= + ( + (unsigned int )(((unsigned int )(IV0 & 0xff000000)) >> 24) | + (unsigned int )((unsigned int )(IV0 & 0x00ff0000) >> 8) | + (unsigned int )((unsigned int )(IV0 & 0x0000ff00) << 8) | + (unsigned int )((unsigned int )(IV0 & 0x000000ff) << 24) + ) ^ hash0; + + schedule1= + ( + (unsigned int )(((unsigned int )(IV4 & 0xff000000)) >> 24) | + (unsigned int )((unsigned int )(IV4 & 0x00ff0000) >> 8) | + (unsigned int )((unsigned int )(IV4 & 0x0000ff00) << 8) | + (unsigned int )((unsigned int )(IV4 & 0x000000ff) << 24) + ) ^ hash1; + + schedule2= + ( + (unsigned int )(((unsigned int )(IV8 & 0xff000000)) >> 24) | + (unsigned int )((unsigned int )(IV8 & 0x00ff0000) >> 8) | + (unsigned int )((unsigned int )(IV8 & 0x0000ff00) << 8) | + (unsigned int )((unsigned int )(IV8 & 0x000000ff) << 24) + ) ^ hash2; + + schedule3= + ( + (unsigned int )(((unsigned int )(IV12 & 0xff000000)) >> 24) | + (unsigned int )((unsigned int )(IV12 & 0x00ff0000) >> 8) | + (unsigned int )((unsigned int )(IV12 & 0x0000ff00) << 8) | + (unsigned int )((unsigned int )(IV12 & 0x000000ff) << 24) + ) ^ hash3; + schedule4 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule0 >> 24], TS1[(schedule1 >> 16) & 0xFF], TS2[(schedule2 >> 8) & 0xFF]) , TS3[schedule3 & 0xFF] , hash4); schedule5 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule1 >> 24], TS1[(schedule2 >> 16) & 0xFF], TS2[(schedule3 >> 8) & 0xFF]) , TS3[schedule0 & 0xFF] , hash5); schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6); @@ -1397,3 +1435,1113 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * return; } + +void encrypt( + unsigned int k0, unsigned int k1, unsigned int k2, unsigned int k3, unsigned int k4, unsigned int k5, unsigned int k6, unsigned int k7, + unsigned int m0, unsigned int m1, unsigned int m2, unsigned int m3, + unsigned int * output0, unsigned int * output1, unsigned int * output2, unsigned int * output3 +) +{ + unsigned int enc_schedule0, enc_schedule1, enc_schedule2, enc_schedule3, enc_schedule4, enc_schedule5, enc_schedule6, enc_schedule7; + unsigned int local_key0, local_key1, local_key2, local_key3, local_key4, local_key5, local_key6, local_key7; + + local_key0=k0; + local_key1=k1; + local_key2=k2; + local_key3=k3; + local_key4=k4; + local_key5=k5; + local_key6=k6; + local_key7=k7; + + enc_schedule0=(unsigned int )(((unsigned int )(m0 & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(m0 & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(m0 & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(m0 & 0x000000ff) << 24); + enc_schedule0 = enc_schedule0 ^ local_key0; + + enc_schedule1=(unsigned int )(((unsigned int )(m1 & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(m1 & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(m1 & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(m1 & 0x000000ff) << 24); + enc_schedule1 = enc_schedule1 ^ local_key1; + + enc_schedule2=(unsigned int )(((unsigned int )(m2 & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(m2 & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(m2 & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(m2 & 0x000000ff) << 24); + enc_schedule2 = enc_schedule2 ^ local_key2; + + enc_schedule3=(unsigned int )(((unsigned int )(m3 & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(m3 & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(m3 & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(m3 & 0x000000ff) << 24); + enc_schedule3 = enc_schedule3 ^ local_key3; + + enc_schedule4 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule0 >> 24], TS1[(enc_schedule1 >> 16) & 0xFF], TS2[(enc_schedule2 >> 8) & 0xFF]) , TS3[enc_schedule3 & 0xFF] , local_key4); + enc_schedule5 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule1 >> 24], TS1[(enc_schedule2 >> 16) & 0xFF], TS2[(enc_schedule3 >> 8) & 0xFF]) , TS3[enc_schedule0 & 0xFF] , local_key5); + enc_schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule2 >> 24], TS1[(enc_schedule3 >> 16) & 0xFF], TS2[(enc_schedule0 >> 8) & 0xFF]) , TS3[enc_schedule1 & 0xFF] , local_key6); + enc_schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule3 >> 24], TS1[(enc_schedule0 >> 16) & 0xFF], TS2[(enc_schedule1 >> 8) & 0xFF]) , TS3[enc_schedule2 & 0xFF] , local_key7); + + local_key0 ^= LOP3LUT_XOR( + LOP3LUT_XOR( (TS2[(local_key7 >> 24) ] & 0x000000FF), (TS3[(local_key7 >> 16) & 0xFF] & 0xFF000000), (TS0[(local_key7 >> 8) & 0xFF] & 0x00FF0000)), + (TS1[(local_key7 ) & 0xFF] & 0x0000FF00), 0x01000000 + ); //RCON[0]; + local_key1 ^= local_key0; local_key2 ^= local_key1; local_key3 ^= local_key2; + + enc_schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule4 >> 24], TS1[(enc_schedule5 >> 16) & 0xFF], TS2[(enc_schedule6 >> 8) & 0xFF]) , TS3[enc_schedule7 & 0xFF] , local_key0); + enc_schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule5 >> 24], TS1[(enc_schedule6 >> 16) & 0xFF], TS2[(enc_schedule7 >> 8) & 0xFF]) , TS3[enc_schedule4 & 0xFF] , local_key1); + enc_schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule6 >> 24], TS1[(enc_schedule7 >> 16) & 0xFF], TS2[(enc_schedule4 >> 8) & 0xFF]) , TS3[enc_schedule5 & 0xFF] , local_key2); + enc_schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule7 >> 24], TS1[(enc_schedule4 >> 16) & 0xFF], TS2[(enc_schedule5 >> 8) & 0xFF]) , TS3[enc_schedule6 & 0xFF] , local_key3); + + local_key4 ^= (TS3[(local_key3 >> 24) ] & 0xFF000000) ^ + (TS0[(local_key3 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key3 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS2[(local_key3 ) & 0xFF] & 0x000000FF); + local_key5 ^= local_key4; + local_key6 ^= local_key5; + local_key7 ^= local_key6; + + enc_schedule4 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule0 >> 24], TS1[(enc_schedule1 >> 16) & 0xFF], TS2[(enc_schedule2 >> 8) & 0xFF]) , TS3[enc_schedule3 & 0xFF] , local_key4); + enc_schedule5 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule1 >> 24], TS1[(enc_schedule2 >> 16) & 0xFF], TS2[(enc_schedule3 >> 8) & 0xFF]) , TS3[enc_schedule0 & 0xFF] , local_key5); + enc_schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule2 >> 24], TS1[(enc_schedule3 >> 16) & 0xFF], TS2[(enc_schedule0 >> 8) & 0xFF]) , TS3[enc_schedule1 & 0xFF] , local_key6); + enc_schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule3 >> 24], TS1[(enc_schedule0 >> 16) & 0xFF], TS2[(enc_schedule1 >> 8) & 0xFF]) , TS3[enc_schedule2 & 0xFF] , local_key7); + + local_key0 ^= (TS2[(local_key7 >> 24) ] & 0x000000FF) ^ + (TS3[(local_key7 >> 16) & 0xFF] & 0xFF000000) ^ + (TS0[(local_key7 >> 8) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key7 ) & 0xFF] & 0x0000FF00) ^ 0x02000000; //RCON[1]; + local_key1 ^= local_key0; local_key2 ^= local_key1; local_key3 ^= local_key2; + + enc_schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule4 >> 24], TS1[(enc_schedule5 >> 16) & 0xFF], TS2[(enc_schedule6 >> 8) & 0xFF]) , TS3[enc_schedule7 & 0xFF] , local_key0); + enc_schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule5 >> 24], TS1[(enc_schedule6 >> 16) & 0xFF], TS2[(enc_schedule7 >> 8) & 0xFF]) , TS3[enc_schedule4 & 0xFF] , local_key1); + enc_schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule6 >> 24], TS1[(enc_schedule7 >> 16) & 0xFF], TS2[(enc_schedule4 >> 8) & 0xFF]) , TS3[enc_schedule5 & 0xFF] , local_key2); + enc_schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule7 >> 24], TS1[(enc_schedule4 >> 16) & 0xFF], TS2[(enc_schedule5 >> 8) & 0xFF]) , TS3[enc_schedule6 & 0xFF] , local_key3); + + local_key4 ^= (TS3[(local_key3 >> 24) ] & 0xFF000000) ^ + (TS0[(local_key3 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key3 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS2[(local_key3 ) & 0xFF] & 0x000000FF); + local_key5 ^= local_key4; + local_key6 ^= local_key5; + local_key7 ^= local_key6; + + enc_schedule4 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule0 >> 24], TS1[(enc_schedule1 >> 16) & 0xFF], TS2[(enc_schedule2 >> 8) & 0xFF]) , TS3[enc_schedule3 & 0xFF] , local_key4); + enc_schedule5 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule1 >> 24], TS1[(enc_schedule2 >> 16) & 0xFF], TS2[(enc_schedule3 >> 8) & 0xFF]) , TS3[enc_schedule0 & 0xFF] , local_key5); + enc_schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule2 >> 24], TS1[(enc_schedule3 >> 16) & 0xFF], TS2[(enc_schedule0 >> 8) & 0xFF]) , TS3[enc_schedule1 & 0xFF] , local_key6); + enc_schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule3 >> 24], TS1[(enc_schedule0 >> 16) & 0xFF], TS2[(enc_schedule1 >> 8) & 0xFF]) , TS3[enc_schedule2 & 0xFF] , local_key7); + + + local_key0 ^= (TS2[(local_key7 >> 24) ] & 0x000000FF) ^ + (TS3[(local_key7 >> 16) & 0xFF] & 0xFF000000) ^ + (TS0[(local_key7 >> 8) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key7 ) & 0xFF] & 0x0000FF00) ^ 0x04000000; //RCON[2]; + local_key1 ^= local_key0; local_key2 ^= local_key1; local_key3 ^= local_key2; + + enc_schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule4 >> 24], TS1[(enc_schedule5 >> 16) & 0xFF], TS2[(enc_schedule6 >> 8) & 0xFF]) , TS3[enc_schedule7 & 0xFF] , local_key0); + enc_schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule5 >> 24], TS1[(enc_schedule6 >> 16) & 0xFF], TS2[(enc_schedule7 >> 8) & 0xFF]) , TS3[enc_schedule4 & 0xFF] , local_key1); + enc_schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule6 >> 24], TS1[(enc_schedule7 >> 16) & 0xFF], TS2[(enc_schedule4 >> 8) & 0xFF]) , TS3[enc_schedule5 & 0xFF] , local_key2); + enc_schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule7 >> 24], TS1[(enc_schedule4 >> 16) & 0xFF], TS2[(enc_schedule5 >> 8) & 0xFF]) , TS3[enc_schedule6 & 0xFF] , local_key3); + + + local_key4 ^= (TS3[(local_key3 >> 24) ] & 0xFF000000) ^ + (TS0[(local_key3 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key3 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS2[(local_key3 ) & 0xFF] & 0x000000FF); + local_key5 ^= local_key4; + local_key6 ^= local_key5; + local_key7 ^= local_key6; + + enc_schedule4 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule0 >> 24], TS1[(enc_schedule1 >> 16) & 0xFF], TS2[(enc_schedule2 >> 8) & 0xFF]) , TS3[enc_schedule3 & 0xFF] , local_key4); + enc_schedule5 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule1 >> 24], TS1[(enc_schedule2 >> 16) & 0xFF], TS2[(enc_schedule3 >> 8) & 0xFF]) , TS3[enc_schedule0 & 0xFF] , local_key5); + enc_schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule2 >> 24], TS1[(enc_schedule3 >> 16) & 0xFF], TS2[(enc_schedule0 >> 8) & 0xFF]) , TS3[enc_schedule1 & 0xFF] , local_key6); + enc_schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule3 >> 24], TS1[(enc_schedule0 >> 16) & 0xFF], TS2[(enc_schedule1 >> 8) & 0xFF]) , TS3[enc_schedule2 & 0xFF] , local_key7); + + local_key0 ^= (TS2[(local_key7 >> 24) ] & 0x000000FF) ^ + (TS3[(local_key7 >> 16) & 0xFF] & 0xFF000000) ^ + (TS0[(local_key7 >> 8) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key7 ) & 0xFF] & 0x0000FF00) ^ 0x08000000; //RCON[3]; + local_key1 ^= local_key0; local_key2 ^= local_key1; local_key3 ^= local_key2; + + enc_schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule4 >> 24], TS1[(enc_schedule5 >> 16) & 0xFF], TS2[(enc_schedule6 >> 8) & 0xFF]) , TS3[enc_schedule7 & 0xFF] , local_key0); + enc_schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule5 >> 24], TS1[(enc_schedule6 >> 16) & 0xFF], TS2[(enc_schedule7 >> 8) & 0xFF]) , TS3[enc_schedule4 & 0xFF] , local_key1); + enc_schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule6 >> 24], TS1[(enc_schedule7 >> 16) & 0xFF], TS2[(enc_schedule4 >> 8) & 0xFF]) , TS3[enc_schedule5 & 0xFF] , local_key2); + enc_schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule7 >> 24], TS1[(enc_schedule4 >> 16) & 0xFF], TS2[(enc_schedule5 >> 8) & 0xFF]) , TS3[enc_schedule6 & 0xFF] , local_key3); + + local_key4 ^= (TS3[(local_key3 >> 24) ] & 0xFF000000) ^ + (TS0[(local_key3 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key3 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS2[(local_key3 ) & 0xFF] & 0x000000FF); + local_key5 ^= local_key4; + local_key6 ^= local_key5; + local_key7 ^= local_key6; + + enc_schedule4 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule0 >> 24], TS1[(enc_schedule1 >> 16) & 0xFF], TS2[(enc_schedule2 >> 8) & 0xFF]) , TS3[enc_schedule3 & 0xFF] , local_key4); + enc_schedule5 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule1 >> 24], TS1[(enc_schedule2 >> 16) & 0xFF], TS2[(enc_schedule3 >> 8) & 0xFF]) , TS3[enc_schedule0 & 0xFF] , local_key5); + enc_schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule2 >> 24], TS1[(enc_schedule3 >> 16) & 0xFF], TS2[(enc_schedule0 >> 8) & 0xFF]) , TS3[enc_schedule1 & 0xFF] , local_key6); + enc_schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule3 >> 24], TS1[(enc_schedule0 >> 16) & 0xFF], TS2[(enc_schedule1 >> 8) & 0xFF]) , TS3[enc_schedule2 & 0xFF] , local_key7); + + local_key0 ^= (TS2[(local_key7 >> 24) ] & 0x000000FF) ^ + (TS3[(local_key7 >> 16) & 0xFF] & 0xFF000000) ^ + (TS0[(local_key7 >> 8) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key7 ) & 0xFF] & 0x0000FF00) ^ 0x10000000; //RCON[4]; + local_key1 ^= local_key0; local_key2 ^= local_key1; local_key3 ^= local_key2; + + enc_schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule4 >> 24], TS1[(enc_schedule5 >> 16) & 0xFF], TS2[(enc_schedule6 >> 8) & 0xFF]) , TS3[enc_schedule7 & 0xFF] , local_key0); + enc_schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule5 >> 24], TS1[(enc_schedule6 >> 16) & 0xFF], TS2[(enc_schedule7 >> 8) & 0xFF]) , TS3[enc_schedule4 & 0xFF] , local_key1); + enc_schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule6 >> 24], TS1[(enc_schedule7 >> 16) & 0xFF], TS2[(enc_schedule4 >> 8) & 0xFF]) , TS3[enc_schedule5 & 0xFF] , local_key2); + enc_schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule7 >> 24], TS1[(enc_schedule4 >> 16) & 0xFF], TS2[(enc_schedule5 >> 8) & 0xFF]) , TS3[enc_schedule6 & 0xFF] , local_key3); + + local_key4 ^= (TS3[(local_key3 >> 24) ] & 0xFF000000) ^ + (TS0[(local_key3 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key3 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS2[(local_key3 ) & 0xFF] & 0x000000FF); + local_key5 ^= local_key4; + local_key6 ^= local_key5; + local_key7 ^= local_key6; + + enc_schedule4 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule0 >> 24], TS1[(enc_schedule1 >> 16) & 0xFF], TS2[(enc_schedule2 >> 8) & 0xFF]) , TS3[enc_schedule3 & 0xFF] , local_key4); + enc_schedule5 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule1 >> 24], TS1[(enc_schedule2 >> 16) & 0xFF], TS2[(enc_schedule3 >> 8) & 0xFF]) , TS3[enc_schedule0 & 0xFF] , local_key5); + enc_schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule2 >> 24], TS1[(enc_schedule3 >> 16) & 0xFF], TS2[(enc_schedule0 >> 8) & 0xFF]) , TS3[enc_schedule1 & 0xFF] , local_key6); + enc_schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule3 >> 24], TS1[(enc_schedule0 >> 16) & 0xFF], TS2[(enc_schedule1 >> 8) & 0xFF]) , TS3[enc_schedule2 & 0xFF] , local_key7); + + + local_key0 ^= (TS2[(local_key7 >> 24) ] & 0x000000FF) ^ + (TS3[(local_key7 >> 16) & 0xFF] & 0xFF000000) ^ + (TS0[(local_key7 >> 8) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key7 ) & 0xFF] & 0x0000FF00) ^ 0x20000000; //RCON[5]; + local_key1 ^= local_key0; local_key2 ^= local_key1; local_key3 ^= local_key2; + + enc_schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule4 >> 24], TS1[(enc_schedule5 >> 16) & 0xFF], TS2[(enc_schedule6 >> 8) & 0xFF]) , TS3[enc_schedule7 & 0xFF] , local_key0); + enc_schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule5 >> 24], TS1[(enc_schedule6 >> 16) & 0xFF], TS2[(enc_schedule7 >> 8) & 0xFF]) , TS3[enc_schedule4 & 0xFF] , local_key1); + enc_schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule6 >> 24], TS1[(enc_schedule7 >> 16) & 0xFF], TS2[(enc_schedule4 >> 8) & 0xFF]) , TS3[enc_schedule5 & 0xFF] , local_key2); + enc_schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule7 >> 24], TS1[(enc_schedule4 >> 16) & 0xFF], TS2[(enc_schedule5 >> 8) & 0xFF]) , TS3[enc_schedule6 & 0xFF] , local_key3); + + local_key4 ^= (TS3[(local_key3 >> 24)] & 0xFF000000) ^ + (TS0[(local_key3 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key3 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS2[(local_key3 ) & 0xFF] & 0x000000FF); + local_key5 ^= local_key4; + local_key6 ^= local_key5; + local_key7 ^= local_key6; + + enc_schedule4 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule0 >> 24], TS1[(enc_schedule1 >> 16) & 0xFF], TS2[(enc_schedule2 >> 8) & 0xFF]) , TS3[enc_schedule3 & 0xFF] , local_key4); + enc_schedule5 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule1 >> 24], TS1[(enc_schedule2 >> 16) & 0xFF], TS2[(enc_schedule3 >> 8) & 0xFF]) , TS3[enc_schedule0 & 0xFF] , local_key5); + enc_schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule2 >> 24], TS1[(enc_schedule3 >> 16) & 0xFF], TS2[(enc_schedule0 >> 8) & 0xFF]) , TS3[enc_schedule1 & 0xFF] , local_key6); + enc_schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule3 >> 24], TS1[(enc_schedule0 >> 16) & 0xFF], TS2[(enc_schedule1 >> 8) & 0xFF]) , TS3[enc_schedule2 & 0xFF] , local_key7); + + local_key0 ^= (TS2[(local_key7 >> 24)] & 0x000000FF) ^ + (TS3[(local_key7 >> 16) & 0xFF] & 0xFF000000) ^ + (TS0[(local_key7 >> 8) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key7 ) & 0xFF] & 0x0000FF00) ^ 0x40000000; //RCON[6]; + local_key1 ^= local_key0; + local_key2 ^= local_key1; + local_key3 ^= local_key2; + + enc_schedule0 = (TS2[(enc_schedule4 >> 24) ] & 0xFF000000) ^ + (TS3[(enc_schedule5 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS0[(enc_schedule6 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS1[(enc_schedule7 ) & 0xFF] & 0x000000FF) ^ local_key0; + + enc_schedule1 = (TS2[(enc_schedule5 >> 24) ] & 0xFF000000) ^ + (TS3[(enc_schedule6 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS0[(enc_schedule7 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS1[(enc_schedule4 ) & 0xFF] & 0x000000FF) ^ local_key1; + + enc_schedule2 = (TS2[(enc_schedule6 >> 24) ] & 0xFF000000) ^ + (TS3[(enc_schedule7 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS0[(enc_schedule4 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS1[(enc_schedule5 ) & 0xFF] & 0x000000FF) ^ local_key2; + + enc_schedule3 = (TS2[(enc_schedule7 >> 24) ] & 0xFF000000) ^ + (TS3[(enc_schedule4 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS0[(enc_schedule5 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS1[(enc_schedule6 ) & 0xFF] & 0x000000FF) ^ local_key3; + + output0[0]=(unsigned int )(((unsigned int )(enc_schedule0 & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(enc_schedule0 & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(enc_schedule0 & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(enc_schedule0 & 0x000000ff) << 24); + output1[0]=(unsigned int )(((unsigned int )(enc_schedule1 & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(enc_schedule1 & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(enc_schedule1 & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(enc_schedule1 & 0x000000ff) << 24); + output2[0]=(unsigned int )(((unsigned int )(enc_schedule2 & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(enc_schedule2 & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(enc_schedule2 & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(enc_schedule2 & 0x000000ff) << 24); + output3[0]=(unsigned int )(((unsigned int )(enc_schedule3 & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(enc_schedule3 & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(enc_schedule3 & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(enc_schedule3 & 0x000000ff) << 24); +} + +__kernel void opencl_bitcracker_attack_mac(int numPassword, __global unsigned char *w_password, + __global int *found, __global unsigned char * vmkKey, + __global unsigned int *w_blocks_d, + unsigned int IV0, unsigned int IV4, unsigned int IV8, unsigned int IV12, + __global unsigned char * mac, + unsigned int macIV0, unsigned int macIV4, unsigned int macIV8, unsigned int macIV12, + unsigned int computeMacIV0, unsigned int computeMacIV4, unsigned int computeMacIV8, unsigned int computeMacIV12 + ) +{ + unsigned int schedule0, schedule1, schedule2, schedule3, schedule4, schedule5, schedule6, schedule7, schedule8, schedule9; + unsigned int schedule10, schedule11, schedule12, schedule13, schedule14, schedule15, schedule16, schedule17, schedule18, schedule19; + unsigned int schedule20, schedule21, schedule22, schedule23, schedule24, schedule25, schedule26, schedule27, schedule28, schedule29; + unsigned int schedule30, schedule31; + unsigned int first_hash0, first_hash1, first_hash2, first_hash3, first_hash4, first_hash5, first_hash6, first_hash7; + unsigned int hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7; + unsigned int a, b, c, d, e, f, g, h; + + int gIndex = (int)get_global_id(0); + int index_generic; + int indexW=(gIndex*FIXED_PASSWORD_BUFFER); + int curr_fetch=0; + + while(gIndex < numPassword) + { + first_hash0 = 0x6A09E667; + first_hash1 = 0xBB67AE85; + first_hash2 = 0x3C6EF372; + first_hash3 = 0xA54FF53A; + first_hash4 = 0x510E527F; + first_hash5 = 0x9B05688C; + first_hash6 = 0x1F83D9AB; + first_hash7 = 0x5BE0CD19; + + //----------------------------------------------------- FIRST HASH ------------------------------------------------ + a = 0x6A09E667; + b = 0xBB67AE85; + c = 0x3C6EF372; + d = 0xA54FF53A; + e = 0x510E527F; + f = 0x9B05688C; + g = 0x1F83D9AB; + h = 0x5BE0CD19; + + indexW=(gIndex*FIXED_PASSWORD_BUFFER); + curr_fetch=0; + index_generic=MAX_INPUT_PASSWORD_LEN; + + //--------------------- SCHEDULE ------------------- + schedule0 = ((unsigned int)w_password[(indexW+curr_fetch)] << 24) | 0 | ((unsigned int)w_password[(indexW+curr_fetch+1)] << 8) | 0; + curr_fetch+=2; + schedule1 = ((unsigned int)w_password[(indexW+curr_fetch)] << 24) | 0 | ((unsigned int)w_password[(indexW+curr_fetch+1)] << 8) | 0; + curr_fetch+=2; + schedule2 = ((unsigned int)w_password[(indexW+curr_fetch)] << 24) | 0 | ((unsigned int)w_password[(indexW+curr_fetch+1)] << 8) | 0; + curr_fetch+=2; + schedule3 = ((unsigned int)w_password[(indexW+curr_fetch)] << 24) | 0 | ((unsigned int)w_password[(indexW+curr_fetch+1)] << 8) | 0; + curr_fetch+=2; + + schedule4 = ((unsigned int)w_password[(indexW+curr_fetch)] << 24) | 0 | ((unsigned int)w_password[(indexW+curr_fetch+1)] << 8) | 0; + if((unsigned int)w_password[(indexW+curr_fetch)] == END_STRING) { index_generic=curr_fetch; } + if((unsigned int)w_password[(indexW+curr_fetch+1)] == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule5 = ((unsigned int)w_password[(indexW+curr_fetch)] << 24) | 0 | ((unsigned int)w_password[(indexW+curr_fetch+1)] << 8) | 0; + if((unsigned int)w_password[(indexW+curr_fetch)] == END_STRING) { index_generic=curr_fetch; } + if((unsigned int)w_password[(indexW+curr_fetch+1)] == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule6 = ((unsigned int)w_password[(indexW+curr_fetch)] << 24) | 0 | ((unsigned int)w_password[(indexW+curr_fetch+1)] << 8) | 0; + if((unsigned int)w_password[(indexW+curr_fetch)] == END_STRING) { index_generic=curr_fetch; } + if((unsigned int)w_password[(indexW+curr_fetch+1)] == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule7 = ((unsigned int)w_password[(indexW+curr_fetch)] << 24) | 0 | ((unsigned int)w_password[(indexW+curr_fetch+1)] << 8) | 0; + if((unsigned int)w_password[(indexW+curr_fetch)] == END_STRING) { index_generic=curr_fetch; } + if((unsigned int)w_password[(indexW+curr_fetch+1)] == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule8 = ((unsigned int)w_password[(indexW+curr_fetch)] << 24) | 0 | ((unsigned int)w_password[(indexW+curr_fetch+1)] << 8) | 0; + if((unsigned int)w_password[(indexW+curr_fetch)] == END_STRING) { index_generic=curr_fetch; } + if((unsigned int)w_password[(indexW+curr_fetch+1)] == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule9 = ((unsigned int)w_password[(indexW+curr_fetch)] << 24) | 0 | ((unsigned int)w_password[(indexW+curr_fetch+1)] << 8) | 0; + if((unsigned int)w_password[(indexW+curr_fetch)] == END_STRING) { index_generic=curr_fetch; } + if((unsigned int)w_password[(indexW+curr_fetch+1)] == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule10 = ((unsigned int)w_password[(indexW+curr_fetch)] << 24) | 0 | ((unsigned int)w_password[(indexW+curr_fetch+1)] << 8) | 0; + if((unsigned int)w_password[(indexW+curr_fetch)] == END_STRING) { index_generic=curr_fetch; } + if((unsigned int)w_password[(indexW+curr_fetch+1)] == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule11 = ((unsigned int)w_password[(indexW+curr_fetch)] << 24) | 0 | ((unsigned int)w_password[(indexW+curr_fetch+1)] << 8) | 0; + if((unsigned int)w_password[(indexW+curr_fetch)] == END_STRING) { index_generic=curr_fetch; } + if((unsigned int)w_password[(indexW+curr_fetch+1)] == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule12 = ((unsigned int)w_password[(indexW+curr_fetch)] << 24) | 0 | ((unsigned int)w_password[(indexW+curr_fetch+1)] << 8) | 0; + if((unsigned int)w_password[(indexW+curr_fetch)] == END_STRING) { index_generic=curr_fetch; } + if((unsigned int)w_password[(indexW+curr_fetch+1)] == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule13 = ((unsigned int)w_password[(indexW+curr_fetch)] << 24) | 0 | ((unsigned int)w_password[(indexW+curr_fetch+1)] << 8) | 0; + if((unsigned int)w_password[(indexW+curr_fetch)] == END_STRING) { index_generic=curr_fetch; } + if((unsigned int)w_password[(indexW+curr_fetch+1)] == END_STRING) { index_generic=curr_fetch+1; } + + if(index_generic == MAX_INPUT_PASSWORD_LEN) schedule13 = schedule13 | ((unsigned int)0x8000); + schedule14=0; + index_generic*=2; + schedule15 = ((unsigned char)((index_generic << 3) >> 8)) << 8 | ((unsigned char)(index_generic << 3)); + + //----------------------------------------------- + + ALL_SCHEDULE_LAST16() + + ROUND(a, b, c, d, e, f, g, h, schedule0, 0x428A2F98) + ROUND(h, a, b, c, d, e, f, g, schedule1, 0x71374491) + ROUND(g, h, a, b, c, d, e, f, schedule2, 0xB5C0FBCF) + ROUND(f, g, h, a, b, c, d, e, schedule3, 0xE9B5DBA5) + ROUND(e, f, g, h, a, b, c, d, schedule4, 0x3956C25B) + ROUND(d, e, f, g, h, a, b, c, schedule5, 0x59F111F1) + ROUND(c, d, e, f, g, h, a, b, schedule6, 0x923F82A4) + ROUND(b, c, d, e, f, g, h, a, schedule7, 0xAB1C5ED5) + ROUND(a, b, c, d, e, f, g, h, schedule8, 0xD807AA98) + ROUND(h, a, b, c, d, e, f, g, schedule9, 0x12835B01) + ROUND(g, h, a, b, c, d, e, f, schedule10, 0x243185BE) + ROUND(f, g, h, a, b, c, d, e, schedule11, 0x550C7DC3) + ROUND(e, f, g, h, a, b, c, d, schedule12, 0x72BE5D74) + ROUND(d, e, f, g, h, a, b, c, schedule13, 0x80DEB1FE) + ROUND(c, d, e, f, g, h, a, b, schedule14, 0x9BDC06A7) + ROUND(b, c, d, e, f, g, h, a, schedule15, 0xC19BF174) + ROUND(a, b, c, d, e, f, g, h, schedule16, 0xE49B69C1) + ROUND(h, a, b, c, d, e, f, g, schedule17, 0xEFBE4786) + ROUND(g, h, a, b, c, d, e, f, schedule18, 0x0FC19DC6) + ROUND(f, g, h, a, b, c, d, e, schedule19, 0x240CA1CC) + ROUND(e, f, g, h, a, b, c, d, schedule20, 0x2DE92C6F) + ROUND(d, e, f, g, h, a, b, c, schedule21, 0x4A7484AA) + ROUND(c, d, e, f, g, h, a, b, schedule22, 0x5CB0A9DC) + ROUND(b, c, d, e, f, g, h, a, schedule23, 0x76F988DA) + ROUND(a, b, c, d, e, f, g, h, schedule24, 0x983E5152) + ROUND(h, a, b, c, d, e, f, g, schedule25, 0xA831C66D) + ROUND(g, h, a, b, c, d, e, f, schedule26, 0xB00327C8) + ROUND(f, g, h, a, b, c, d, e, schedule27, 0xBF597FC7) + ROUND(e, f, g, h, a, b, c, d, schedule28, 0xC6E00BF3) + ROUND(d, e, f, g, h, a, b, c, schedule29, 0xD5A79147) + ROUND(c, d, e, f, g, h, a, b, schedule30, 0x06CA6351) + ROUND(b, c, d, e, f, g, h, a, schedule31, 0x14292967) + + ALL_SCHEDULE32() + + ROUND(a, b, c, d, e, f, g, h, schedule0, 0x27B70A85) + ROUND(h, a, b, c, d, e, f, g, schedule1, 0x2E1B2138) + ROUND(g, h, a, b, c, d, e, f, schedule2, 0x4D2C6DFC) + ROUND(f, g, h, a, b, c, d, e, schedule3, 0x53380D13) + ROUND(e, f, g, h, a, b, c, d, schedule4, 0x650A7354) + ROUND(d, e, f, g, h, a, b, c, schedule5, 0x766A0ABB) + ROUND(c, d, e, f, g, h, a, b, schedule6, 0x81C2C92E) + ROUND(b, c, d, e, f, g, h, a, schedule7, 0x92722C85) + ROUND(a, b, c, d, e, f, g, h, schedule8, 0xA2BFE8A1) + ROUND(h, a, b, c, d, e, f, g, schedule9, 0xA81A664B) + ROUND(g, h, a, b, c, d, e, f, schedule10, 0xC24B8B70) + ROUND(f, g, h, a, b, c, d, e, schedule11, 0xC76C51A3) + ROUND(e, f, g, h, a, b, c, d, schedule12, 0xD192E819) + ROUND(d, e, f, g, h, a, b, c, schedule13, 0xD6990624) + ROUND(c, d, e, f, g, h, a, b, schedule14, 0xF40E3585) + ROUND(b, c, d, e, f, g, h, a, schedule15, 0x106AA070) + ROUND(a, b, c, d, e, f, g, h, schedule16, 0x19A4C116) + ROUND(h, a, b, c, d, e, f, g, schedule17, 0x1E376C08) + ROUND(g, h, a, b, c, d, e, f, schedule18, 0x2748774C) + ROUND(f, g, h, a, b, c, d, e, schedule19, 0x34B0BCB5) + ROUND(e, f, g, h, a, b, c, d, schedule20, 0x391C0CB3) + ROUND(d, e, f, g, h, a, b, c, schedule21, 0x4ED8AA4A) + ROUND(c, d, e, f, g, h, a, b, schedule22, 0x5B9CCA4F) + ROUND(b, c, d, e, f, g, h, a, schedule23, 0x682E6FF3) + ROUND(a, b, c, d, e, f, g, h, schedule24, 0x748F82EE) + ROUND(h, a, b, c, d, e, f, g, schedule25, 0x78A5636F) + ROUND(g, h, a, b, c, d, e, f, schedule26, 0x84C87814) + ROUND(f, g, h, a, b, c, d, e, schedule27, 0x8CC70208) + ROUND(e, f, g, h, a, b, c, d, schedule28, 0x90BEFFFA) + ROUND(d, e, f, g, h, a, b, c, schedule29, 0xA4506CEB) + ROUND(c, d, e, f, g, h, a, b, schedule30, 0xBEF9A3F7) + ROUND(b, c, d, e, f, g, h, a, schedule31, 0xC67178F2) + + first_hash0 += a; + first_hash1 += b; + first_hash2 += c; + first_hash3 += d; + first_hash4 += e; + first_hash5 += f; + first_hash6 += g; + first_hash7 += h; + +//----------------------------------------------------- SECOND HASH ------------------------------------------------ + schedule0 = first_hash0; + schedule1 = first_hash1; + schedule2 = first_hash2; + schedule3 = first_hash3; + schedule4 = first_hash4; + schedule5 = first_hash5; + schedule6 = first_hash6; + schedule7 = first_hash7; + schedule8 = 0x80000000; + schedule9 = 0; + schedule10 = 0; + schedule11 = 0; + schedule12 = 0; + schedule13 = 0; + schedule14 = 0; + schedule15 = 0x100; + + first_hash0 = 0x6A09E667; + first_hash1 = 0xBB67AE85; + first_hash2 = 0x3C6EF372; + first_hash3 = 0xA54FF53A; + first_hash4 = 0x510E527F; + first_hash5 = 0x9B05688C; + first_hash6 = 0x1F83D9AB; + first_hash7 = 0x5BE0CD19; + + a = first_hash0; + b = first_hash1; + c = first_hash2; + d = first_hash3; + e = first_hash4; + f = first_hash5; + g = first_hash6; + h = first_hash7; + + ALL_SCHEDULE_LAST16() + + ROUND(a, b, c, d, e, f, g, h, schedule0, 0x428A2F98) + ROUND(h, a, b, c, d, e, f, g, schedule1, 0x71374491) + ROUND(g, h, a, b, c, d, e, f, schedule2, 0xB5C0FBCF) + ROUND(f, g, h, a, b, c, d, e, schedule3, 0xE9B5DBA5) + ROUND(e, f, g, h, a, b, c, d, schedule4, 0x3956C25B) + ROUND(d, e, f, g, h, a, b, c, schedule5, 0x59F111F1) + ROUND(c, d, e, f, g, h, a, b, schedule6, 0x923F82A4) + ROUND(b, c, d, e, f, g, h, a, schedule7, 0xAB1C5ED5) + ROUND(a, b, c, d, e, f, g, h, schedule8, 0xD807AA98) + ROUND(h, a, b, c, d, e, f, g, schedule9, 0x12835B01) + ROUND(g, h, a, b, c, d, e, f, schedule10, 0x243185BE) + ROUND(f, g, h, a, b, c, d, e, schedule11, 0x550C7DC3) + ROUND(e, f, g, h, a, b, c, d, schedule12, 0x72BE5D74) + ROUND(d, e, f, g, h, a, b, c, schedule13, 0x80DEB1FE) + ROUND(c, d, e, f, g, h, a, b, schedule14, 0x9BDC06A7) + ROUND(b, c, d, e, f, g, h, a, schedule15, 0xC19BF174) + ROUND(a, b, c, d, e, f, g, h, schedule16, 0xE49B69C1) + ROUND(h, a, b, c, d, e, f, g, schedule17, 0xEFBE4786) + ROUND(g, h, a, b, c, d, e, f, schedule18, 0x0FC19DC6) + ROUND(f, g, h, a, b, c, d, e, schedule19, 0x240CA1CC) + ROUND(e, f, g, h, a, b, c, d, schedule20, 0x2DE92C6F) + ROUND(d, e, f, g, h, a, b, c, schedule21, 0x4A7484AA) + ROUND(c, d, e, f, g, h, a, b, schedule22, 0x5CB0A9DC) + ROUND(b, c, d, e, f, g, h, a, schedule23, 0x76F988DA) + ROUND(a, b, c, d, e, f, g, h, schedule24, 0x983E5152) + ROUND(h, a, b, c, d, e, f, g, schedule25, 0xA831C66D) + ROUND(g, h, a, b, c, d, e, f, schedule26, 0xB00327C8) + ROUND(f, g, h, a, b, c, d, e, schedule27, 0xBF597FC7) + ROUND(e, f, g, h, a, b, c, d, schedule28, 0xC6E00BF3) + ROUND(d, e, f, g, h, a, b, c, schedule29, 0xD5A79147) + ROUND(c, d, e, f, g, h, a, b, schedule30, 0x06CA6351) + ROUND(b, c, d, e, f, g, h, a, schedule31, 0x14292967) + + ALL_SCHEDULE32() + + ROUND(a, b, c, d, e, f, g, h, schedule0, 0x27B70A85) + ROUND(h, a, b, c, d, e, f, g, schedule1, 0x2E1B2138) + ROUND(g, h, a, b, c, d, e, f, schedule2, 0x4D2C6DFC) + ROUND(f, g, h, a, b, c, d, e, schedule3, 0x53380D13) + ROUND(e, f, g, h, a, b, c, d, schedule4, 0x650A7354) + ROUND(d, e, f, g, h, a, b, c, schedule5, 0x766A0ABB) + ROUND(c, d, e, f, g, h, a, b, schedule6, 0x81C2C92E) + ROUND(b, c, d, e, f, g, h, a, schedule7, 0x92722C85) + ROUND(a, b, c, d, e, f, g, h, schedule8, 0xA2BFE8A1) + ROUND(h, a, b, c, d, e, f, g, schedule9, 0xA81A664B) + ROUND(g, h, a, b, c, d, e, f, schedule10, 0xC24B8B70) + ROUND(f, g, h, a, b, c, d, e, schedule11, 0xC76C51A3) + ROUND(e, f, g, h, a, b, c, d, schedule12, 0xD192E819) + ROUND(d, e, f, g, h, a, b, c, schedule13, 0xD6990624) + ROUND(c, d, e, f, g, h, a, b, schedule14, 0xF40E3585) + ROUND(b, c, d, e, f, g, h, a, schedule15, 0x106AA070) + ROUND(a, b, c, d, e, f, g, h, schedule16, 0x19A4C116) + ROUND(h, a, b, c, d, e, f, g, schedule17, 0x1E376C08) + ROUND(g, h, a, b, c, d, e, f, schedule18, 0x2748774C) + ROUND(f, g, h, a, b, c, d, e, schedule19, 0x34B0BCB5) + ROUND(e, f, g, h, a, b, c, d, schedule20, 0x391C0CB3) + ROUND(d, e, f, g, h, a, b, c, schedule21, 0x4ED8AA4A) + ROUND(c, d, e, f, g, h, a, b, schedule22, 0x5B9CCA4F) + ROUND(b, c, d, e, f, g, h, a, schedule23, 0x682E6FF3) + ROUND(a, b, c, d, e, f, g, h, schedule24, 0x748F82EE) + ROUND(h, a, b, c, d, e, f, g, schedule25, 0x78A5636F) + ROUND(g, h, a, b, c, d, e, f, schedule26, 0x84C87814) + ROUND(f, g, h, a, b, c, d, e, schedule27, 0x8CC70208) + ROUND(e, f, g, h, a, b, c, d, schedule28, 0x90BEFFFA) + ROUND(d, e, f, g, h, a, b, c, schedule29, 0xA4506CEB) + ROUND(c, d, e, f, g, h, a, b, schedule30, 0xBEF9A3F7) + ROUND(b, c, d, e, f, g, h, a, schedule31, 0xC67178F2) + + first_hash0 += a; + first_hash1 += b; + first_hash2 += c; + first_hash3 += d; + first_hash4 += e; + first_hash5 += f; + first_hash6 += g; + first_hash7 += h; + +//----------------------------------------------------- LOOP HASH ------------------------------------------------ + + hash0=0; + hash1=0; + hash2=0; + hash3=0; + hash4=0; + hash5=0; + hash6=0; + hash7=0; + + indexW=0; + + for(index_generic=0; index_generic < ITERATION_NUMBER/2; index_generic++) + { + a = 0x6A09E667; + b = 0xBB67AE85; + c = 0x3C6EF372; + d = 0xA54FF53A; + e = 0x510E527F; + f = 0x9B05688C; + g = 0x1F83D9AB; + h = 0x5BE0CD19; + + schedule0 = hash0; + schedule1 = hash1; + schedule2 = hash2; + schedule3 = hash3; + schedule4 = hash4; + schedule5 = hash5; + schedule6 = hash6; + schedule7 = hash7; + + schedule8 = first_hash0; + schedule9 = first_hash1; + schedule10 = first_hash2; + schedule11 = first_hash3; + schedule12 = first_hash4; + schedule13 = first_hash5; + schedule14 = first_hash6; + schedule15 = first_hash7; + + ALL_SCHEDULE_LAST16() + + ROUND(a, b, c, d, e, f, g, h, schedule0, 0x428A2F98) + ROUND(h, a, b, c, d, e, f, g, schedule1, 0x71374491) + ROUND(g, h, a, b, c, d, e, f, schedule2, 0xB5C0FBCF) + ROUND(f, g, h, a, b, c, d, e, schedule3, 0xE9B5DBA5) + ROUND(e, f, g, h, a, b, c, d, schedule4, 0x3956C25B) + ROUND(d, e, f, g, h, a, b, c, schedule5, 0x59F111F1) + ROUND(c, d, e, f, g, h, a, b, schedule6, 0x923F82A4) + ROUND(b, c, d, e, f, g, h, a, schedule7, 0xAB1C5ED5) + ROUND(a, b, c, d, e, f, g, h, schedule8, 0xD807AA98) + ROUND(h, a, b, c, d, e, f, g, schedule9, 0x12835B01) + ROUND(g, h, a, b, c, d, e, f, schedule10, 0x243185BE) + ROUND(f, g, h, a, b, c, d, e, schedule11, 0x550C7DC3) + ROUND(e, f, g, h, a, b, c, d, schedule12, 0x72BE5D74) + ROUND(d, e, f, g, h, a, b, c, schedule13, 0x80DEB1FE) + ROUND(c, d, e, f, g, h, a, b, schedule14, 0x9BDC06A7) + ROUND(b, c, d, e, f, g, h, a, schedule15, 0xC19BF174) + ROUND(a, b, c, d, e, f, g, h, schedule16, 0xE49B69C1) + ROUND(h, a, b, c, d, e, f, g, schedule17, 0xEFBE4786) + ROUND(g, h, a, b, c, d, e, f, schedule18, 0x0FC19DC6) + ROUND(f, g, h, a, b, c, d, e, schedule19, 0x240CA1CC) + ROUND(e, f, g, h, a, b, c, d, schedule20, 0x2DE92C6F) + ROUND(d, e, f, g, h, a, b, c, schedule21, 0x4A7484AA) + ROUND(c, d, e, f, g, h, a, b, schedule22, 0x5CB0A9DC) + ROUND(b, c, d, e, f, g, h, a, schedule23, 0x76F988DA) + ROUND(a, b, c, d, e, f, g, h, schedule24, 0x983E5152) + ROUND(h, a, b, c, d, e, f, g, schedule25, 0xA831C66D) + ROUND(g, h, a, b, c, d, e, f, schedule26, 0xB00327C8) + ROUND(f, g, h, a, b, c, d, e, schedule27, 0xBF597FC7) + ROUND(e, f, g, h, a, b, c, d, schedule28, 0xC6E00BF3) + ROUND(d, e, f, g, h, a, b, c, schedule29, 0xD5A79147) + ROUND(c, d, e, f, g, h, a, b, schedule30, 0x06CA6351) + ROUND(b, c, d, e, f, g, h, a, schedule31, 0x14292967) + + ALL_SCHEDULE32() + + ROUND(a, b, c, d, e, f, g, h, schedule0, 0x27B70A85) + ROUND(h, a, b, c, d, e, f, g, schedule1, 0x2E1B2138) + ROUND(g, h, a, b, c, d, e, f, schedule2, 0x4D2C6DFC) + ROUND(f, g, h, a, b, c, d, e, schedule3, 0x53380D13) + ROUND(e, f, g, h, a, b, c, d, schedule4, 0x650A7354) + ROUND(d, e, f, g, h, a, b, c, schedule5, 0x766A0ABB) + ROUND(c, d, e, f, g, h, a, b, schedule6, 0x81C2C92E) + ROUND(b, c, d, e, f, g, h, a, schedule7, 0x92722C85) + ROUND(a, b, c, d, e, f, g, h, schedule8, 0xA2BFE8A1) + ROUND(h, a, b, c, d, e, f, g, schedule9, 0xA81A664B) + ROUND(g, h, a, b, c, d, e, f, schedule10, 0xC24B8B70) + ROUND(f, g, h, a, b, c, d, e, schedule11, 0xC76C51A3) + ROUND(e, f, g, h, a, b, c, d, schedule12, 0xD192E819) + ROUND(d, e, f, g, h, a, b, c, schedule13, 0xD6990624) + ROUND(c, d, e, f, g, h, a, b, schedule14, 0xF40E3585) + ROUND(b, c, d, e, f, g, h, a, schedule15, 0x106AA070) + ROUND(a, b, c, d, e, f, g, h, schedule16, 0x19A4C116) + ROUND(h, a, b, c, d, e, f, g, schedule17, 0x1E376C08) + ROUND(g, h, a, b, c, d, e, f, schedule18, 0x2748774C) + ROUND(f, g, h, a, b, c, d, e, schedule19, 0x34B0BCB5) + ROUND(e, f, g, h, a, b, c, d, schedule20, 0x391C0CB3) + ROUND(d, e, f, g, h, a, b, c, schedule21, 0x4ED8AA4A) + ROUND(c, d, e, f, g, h, a, b, schedule22, 0x5B9CCA4F) + ROUND(b, c, d, e, f, g, h, a, schedule23, 0x682E6FF3) + ROUND(a, b, c, d, e, f, g, h, schedule24, 0x748F82EE) + ROUND(h, a, b, c, d, e, f, g, schedule25, 0x78A5636F) + ROUND(g, h, a, b, c, d, e, f, schedule26, 0x84C87814) + ROUND(f, g, h, a, b, c, d, e, schedule27, 0x8CC70208) + ROUND(e, f, g, h, a, b, c, d, schedule28, 0x90BEFFFA) + ROUND(d, e, f, g, h, a, b, c, schedule29, 0xA4506CEB) + ROUND(c, d, e, f, g, h, a, b, schedule30, 0xBEF9A3F7) + ROUND(b, c, d, e, f, g, h, a, schedule31, 0xC67178F2) + + hash0 = 0x6A09E667 + a; + hash1 = 0xBB67AE85 + b; + hash2 = 0x3C6EF372 + c; + hash3 = 0xA54FF53A + d; + hash4 = 0x510E527F + e; + hash5 = 0x9B05688C + f; + hash6 = 0x1F83D9AB + g; + hash7 = 0x5BE0CD19 + h; + + a = hash0; + b = hash1; + c = hash2; + d = hash3; + e = hash4; + f = hash5; + g = hash6; + h = hash7; + + ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 0, 0x428A2F98, 0) + ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 1, 0x71374491, 0) + ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 2, 0xB5C0FBCF, 0) + ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 3, 0xE9B5DBA5, 0) + ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 4, 0x3956C25B, indexW) + ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 5, 0x59F111F1, indexW) + ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 6, 0x923F82A4, indexW) + ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 7, 0xAB1C5ED5, indexW) + ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 8, 0xD807AA98, indexW) + ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 9, 0x12835B01, indexW) + ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 10, 0x243185BE, indexW) + ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 11, 0x550C7DC3, indexW) + ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 12, 0x72BE5D74, indexW) + ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 13, 0x80DEB1FE, indexW) + ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 14, 0x9BDC06A7, indexW) + ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 15, 0xC19BF174, indexW) + ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 16, 0xE49B69C1, indexW) + ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 17, 0xEFBE4786, indexW) + ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 18, 0x0FC19DC6, indexW) + ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 19, 0x240CA1CC, indexW) + ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 20, 0x2DE92C6F, indexW) + ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 21, 0x4A7484AA, indexW) + ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 22, 0x5CB0A9DC, indexW) + ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 23, 0x76F988DA, indexW) + ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 24, 0x983E5152, indexW) + ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 25, 0xA831C66D, indexW) + ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 26, 0xB00327C8, indexW) + ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 27, 0xBF597FC7, indexW) + ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 28, 0xC6E00BF3, indexW) + ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 29, 0xD5A79147, indexW) + ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 30, 0x06CA6351, indexW) + ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 31, 0x14292967, indexW) + + ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 32, 0x27B70A85, indexW) + ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 33, 0x2E1B2138, indexW) + ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 34, 0x4D2C6DFC, indexW) + ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 35, 0x53380D13, indexW) + ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 36, 0x650A7354, indexW) + ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 37, 0x766A0ABB, indexW) + ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 38, 0x81C2C92E, indexW) + ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 39, 0x92722C85, indexW) + ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 40, 0xA2BFE8A1, indexW) + ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 41, 0xA81A664B, indexW) + ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 42, 0xC24B8B70, indexW) + ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 43, 0xC76C51A3, indexW) + ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 44, 0xD192E819, indexW) + ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 45, 0xD6990624, indexW) + ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 46, 0xF40E3585, indexW) + ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 47, 0x106AA070, indexW) + ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 48, 0x19A4C116, indexW) + ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 49, 0x1E376C08, indexW) + ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 50, 0x2748774C, indexW) + ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 51, 0x34B0BCB5, indexW) + ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 52, 0x391C0CB3, indexW) + ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 53, 0x4ED8AA4A, indexW) + ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 54, 0x5B9CCA4F, indexW) + ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 55, 0x682E6FF3, indexW) + ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 56, 0x748F82EE, indexW) + ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 57, 0x78A5636F, indexW) + ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 58, 0x84C87814, indexW) + ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 59, 0x8CC70208, indexW) + ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 60, 0x90BEFFFA, indexW) + ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 61, 0xA4506CEB, indexW) + ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 62, 0xBEF9A3F7, indexW) + ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 63, 0xC67178F2, indexW) + + hash0 += a; + hash1 += b; + hash2 += c; + hash3 += d; + hash4 += e; + hash5 += f; + hash6 += g; + hash7 += h; + + indexW += SINGLE_BLOCK_W_SIZE; + } + + for(index_generic=ITERATION_NUMBER/2; index_generic < ITERATION_NUMBER; index_generic++) + { + a = 0x6A09E667; + b = 0xBB67AE85; + c = 0x3C6EF372; + d = 0xA54FF53A; + e = 0x510E527F; + f = 0x9B05688C; + g = 0x1F83D9AB; + h = 0x5BE0CD19; + + schedule0 = hash0; + schedule1 = hash1; + schedule2 = hash2; + schedule3 = hash3; + schedule4 = hash4; + schedule5 = hash5; + schedule6 = hash6; + schedule7 = hash7; + + schedule8 = first_hash0; + schedule9 = first_hash1; + schedule10 = first_hash2; + schedule11 = first_hash3; + schedule12 = first_hash4; + schedule13 = first_hash5; + schedule14 = first_hash6; + schedule15 = first_hash7; + + ALL_SCHEDULE_LAST16() + + ROUND(a, b, c, d, e, f, g, h, schedule0, 0x428A2F98) + ROUND(h, a, b, c, d, e, f, g, schedule1, 0x71374491) + ROUND(g, h, a, b, c, d, e, f, schedule2, 0xB5C0FBCF) + ROUND(f, g, h, a, b, c, d, e, schedule3, 0xE9B5DBA5) + ROUND(e, f, g, h, a, b, c, d, schedule4, 0x3956C25B) + ROUND(d, e, f, g, h, a, b, c, schedule5, 0x59F111F1) + ROUND(c, d, e, f, g, h, a, b, schedule6, 0x923F82A4) + ROUND(b, c, d, e, f, g, h, a, schedule7, 0xAB1C5ED5) + ROUND(a, b, c, d, e, f, g, h, schedule8, 0xD807AA98) + ROUND(h, a, b, c, d, e, f, g, schedule9, 0x12835B01) + ROUND(g, h, a, b, c, d, e, f, schedule10, 0x243185BE) + ROUND(f, g, h, a, b, c, d, e, schedule11, 0x550C7DC3) + ROUND(e, f, g, h, a, b, c, d, schedule12, 0x72BE5D74) + ROUND(d, e, f, g, h, a, b, c, schedule13, 0x80DEB1FE) + ROUND(c, d, e, f, g, h, a, b, schedule14, 0x9BDC06A7) + ROUND(b, c, d, e, f, g, h, a, schedule15, 0xC19BF174) + ROUND(a, b, c, d, e, f, g, h, schedule16, 0xE49B69C1) + ROUND(h, a, b, c, d, e, f, g, schedule17, 0xEFBE4786) + ROUND(g, h, a, b, c, d, e, f, schedule18, 0x0FC19DC6) + ROUND(f, g, h, a, b, c, d, e, schedule19, 0x240CA1CC) + ROUND(e, f, g, h, a, b, c, d, schedule20, 0x2DE92C6F) + ROUND(d, e, f, g, h, a, b, c, schedule21, 0x4A7484AA) + ROUND(c, d, e, f, g, h, a, b, schedule22, 0x5CB0A9DC) + ROUND(b, c, d, e, f, g, h, a, schedule23, 0x76F988DA) + ROUND(a, b, c, d, e, f, g, h, schedule24, 0x983E5152) + ROUND(h, a, b, c, d, e, f, g, schedule25, 0xA831C66D) + ROUND(g, h, a, b, c, d, e, f, schedule26, 0xB00327C8) + ROUND(f, g, h, a, b, c, d, e, schedule27, 0xBF597FC7) + ROUND(e, f, g, h, a, b, c, d, schedule28, 0xC6E00BF3) + ROUND(d, e, f, g, h, a, b, c, schedule29, 0xD5A79147) + ROUND(c, d, e, f, g, h, a, b, schedule30, 0x06CA6351) + ROUND(b, c, d, e, f, g, h, a, schedule31, 0x14292967) + + ALL_SCHEDULE32() + + ROUND(a, b, c, d, e, f, g, h, schedule0, 0x27B70A85) + ROUND(h, a, b, c, d, e, f, g, schedule1, 0x2E1B2138) + ROUND(g, h, a, b, c, d, e, f, schedule2, 0x4D2C6DFC) + ROUND(f, g, h, a, b, c, d, e, schedule3, 0x53380D13) + ROUND(e, f, g, h, a, b, c, d, schedule4, 0x650A7354) + ROUND(d, e, f, g, h, a, b, c, schedule5, 0x766A0ABB) + ROUND(c, d, e, f, g, h, a, b, schedule6, 0x81C2C92E) + ROUND(b, c, d, e, f, g, h, a, schedule7, 0x92722C85) + ROUND(a, b, c, d, e, f, g, h, schedule8, 0xA2BFE8A1) + ROUND(h, a, b, c, d, e, f, g, schedule9, 0xA81A664B) + ROUND(g, h, a, b, c, d, e, f, schedule10, 0xC24B8B70) + ROUND(f, g, h, a, b, c, d, e, schedule11, 0xC76C51A3) + ROUND(e, f, g, h, a, b, c, d, schedule12, 0xD192E819) + ROUND(d, e, f, g, h, a, b, c, schedule13, 0xD6990624) + ROUND(c, d, e, f, g, h, a, b, schedule14, 0xF40E3585) + ROUND(b, c, d, e, f, g, h, a, schedule15, 0x106AA070) + ROUND(a, b, c, d, e, f, g, h, schedule16, 0x19A4C116) + ROUND(h, a, b, c, d, e, f, g, schedule17, 0x1E376C08) + ROUND(g, h, a, b, c, d, e, f, schedule18, 0x2748774C) + ROUND(f, g, h, a, b, c, d, e, schedule19, 0x34B0BCB5) + ROUND(e, f, g, h, a, b, c, d, schedule20, 0x391C0CB3) + ROUND(d, e, f, g, h, a, b, c, schedule21, 0x4ED8AA4A) + ROUND(c, d, e, f, g, h, a, b, schedule22, 0x5B9CCA4F) + ROUND(b, c, d, e, f, g, h, a, schedule23, 0x682E6FF3) + ROUND(a, b, c, d, e, f, g, h, schedule24, 0x748F82EE) + ROUND(h, a, b, c, d, e, f, g, schedule25, 0x78A5636F) + ROUND(g, h, a, b, c, d, e, f, schedule26, 0x84C87814) + ROUND(f, g, h, a, b, c, d, e, schedule27, 0x8CC70208) + ROUND(e, f, g, h, a, b, c, d, schedule28, 0x90BEFFFA) + ROUND(d, e, f, g, h, a, b, c, schedule29, 0xA4506CEB) + ROUND(c, d, e, f, g, h, a, b, schedule30, 0xBEF9A3F7) + ROUND(b, c, d, e, f, g, h, a, schedule31, 0xC67178F2) + + hash0 = 0x6A09E667 + a; + hash1 = 0xBB67AE85 + b; + hash2 = 0x3C6EF372 + c; + hash3 = 0xA54FF53A + d; + hash4 = 0x510E527F + e; + hash5 = 0x9B05688C + f; + hash6 = 0x1F83D9AB + g; + hash7 = 0x5BE0CD19 + h; + + a = hash0; + b = hash1; + c = hash2; + d = hash3; + e = hash4; + f = hash5; + g = hash6; + h = hash7; + + ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 0, 0x428A2F98, 0) + ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 1, 0x71374491, 0) + ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 2, 0xB5C0FBCF, 0) + ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 3, 0xE9B5DBA5, 0) + ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 4, 0x3956C25B, indexW) + ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 5, 0x59F111F1, indexW) + ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 6, 0x923F82A4, indexW) + ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 7, 0xAB1C5ED5, indexW) + ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 8, 0xD807AA98, indexW) + ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 9, 0x12835B01, indexW) + ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 10, 0x243185BE, indexW) + ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 11, 0x550C7DC3, indexW) + ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 12, 0x72BE5D74, indexW) + ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 13, 0x80DEB1FE, indexW) + ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 14, 0x9BDC06A7, indexW) + ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 15, 0xC19BF174, indexW) + ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 16, 0xE49B69C1, indexW) + ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 17, 0xEFBE4786, indexW) + ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 18, 0x0FC19DC6, indexW) + ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 19, 0x240CA1CC, indexW) + ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 20, 0x2DE92C6F, indexW) + ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 21, 0x4A7484AA, indexW) + ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 22, 0x5CB0A9DC, indexW) + ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 23, 0x76F988DA, indexW) + ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 24, 0x983E5152, indexW) + ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 25, 0xA831C66D, indexW) + ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 26, 0xB00327C8, indexW) + ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 27, 0xBF597FC7, indexW) + ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 28, 0xC6E00BF3, indexW) + ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 29, 0xD5A79147, indexW) + ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 30, 0x06CA6351, indexW) + ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 31, 0x14292967, indexW) + + ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 32, 0x27B70A85, indexW) + ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 33, 0x2E1B2138, indexW) + ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 34, 0x4D2C6DFC, indexW) + ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 35, 0x53380D13, indexW) + ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 36, 0x650A7354, indexW) + ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 37, 0x766A0ABB, indexW) + ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 38, 0x81C2C92E, indexW) + ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 39, 0x92722C85, indexW) + ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 40, 0xA2BFE8A1, indexW) + ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 41, 0xA81A664B, indexW) + ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 42, 0xC24B8B70, indexW) + ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 43, 0xC76C51A3, indexW) + ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 44, 0xD192E819, indexW) + ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 45, 0xD6990624, indexW) + ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 46, 0xF40E3585, indexW) + ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 47, 0x106AA070, indexW) + ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 48, 0x19A4C116, indexW) + ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 49, 0x1E376C08, indexW) + ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 50, 0x2748774C, indexW) + ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 51, 0x34B0BCB5, indexW) + ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 52, 0x391C0CB3, indexW) + ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 53, 0x4ED8AA4A, indexW) + ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 54, 0x5B9CCA4F, indexW) + ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 55, 0x682E6FF3, indexW) + ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 56, 0x748F82EE, indexW) + ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 57, 0x78A5636F, indexW) + ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 58, 0x84C87814, indexW) + ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 59, 0x8CC70208, indexW) + ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 60, 0x90BEFFFA, indexW) + ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 61, 0xA4506CEB, indexW) + ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 62, 0xBEF9A3F7, indexW) + ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 63, 0xC67178F2, indexW) + + hash0 += a; + hash1 += b; + hash2 += c; + hash3 += d; + hash4 += e; + hash5 += f; + hash6 += g; + hash7 += h; + + indexW += SINGLE_BLOCK_W_SIZE; + } + + a = IV0; + b = IV4; + c = IV8; + d = IV12; + + encrypt( + hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7, + a, b, c, d, + &(schedule0), &(schedule1), &(schedule2), &(schedule3) + ); + + schedule0= + (((unsigned int)(vmkKey[3] ^ ((unsigned char) (schedule0 >> 24) ))) << 24) | + (((unsigned int)(vmkKey[2] ^ ((unsigned char) (schedule0 >> 16) ))) << 16) | + (((unsigned int)(vmkKey[1] ^ ((unsigned char) (schedule0 >> 8) ))) << 8) | + (((unsigned int)(vmkKey[0] ^ ((unsigned char) (schedule0)))) << 0); + + schedule1= + (((unsigned int)(vmkKey[7] ^ ((unsigned char) (schedule1 >> 24) ))) << 24) | + (((unsigned int)(vmkKey[6] ^ ((unsigned char) (schedule1 >> 16) ))) << 16) | + (((unsigned int)(vmkKey[5] ^ ((unsigned char) (schedule1 >> 8) ))) << 8) | + (((unsigned int)(vmkKey[4] ^ ((unsigned char) (schedule1)))) << 0); + + schedule2= + (((unsigned int)(vmkKey[11] ^ ((unsigned char) (schedule2 >> 24) ))) << 24) | + (((unsigned int)(vmkKey[10] ^ ((unsigned char) (schedule2 >> 16) ))) << 16) | + (((unsigned int)(vmkKey[9] ^ ((unsigned char) (schedule2 >> 8) ))) << 8) | + (((unsigned int)(vmkKey[8] ^ ((unsigned char) (schedule2)))) << 0); + + schedule3= + (((unsigned int)(vmkKey[15] ^ ((unsigned char) (schedule3 >> 24) ))) << 24) | + (((unsigned int)(vmkKey[14] ^ ((unsigned char) (schedule3 >> 16) ))) << 16) | + (((unsigned int)(vmkKey[13] ^ ((unsigned char) (schedule3 >> 8) ))) << 8) | + (((unsigned int)(vmkKey[12] ^ ((unsigned char) (schedule3)))) << 0); + + d += 0x01000000; + + encrypt( + hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7, + a, b, c, d, + &(schedule4), &(schedule5), &(schedule6), &(schedule7) + ); + + schedule4= + (((unsigned int)(vmkKey[19] ^ ((unsigned char) (schedule4 >> 24) ))) << 24) | + (((unsigned int)(vmkKey[18] ^ ((unsigned char) (schedule4 >> 16) ))) << 16) | + (((unsigned int)(vmkKey[17] ^ ((unsigned char) (schedule4 >> 8) ))) << 8) | + (((unsigned int)(vmkKey[16] ^ ((unsigned char) (schedule4)))) << 0); + + schedule5= + (((unsigned int)(vmkKey[23] ^ ((unsigned char) (schedule5 >> 24) ))) << 24) | + (((unsigned int)(vmkKey[22] ^ ((unsigned char) (schedule5 >> 16) ))) << 16) | + (((unsigned int)(vmkKey[21] ^ ((unsigned char) (schedule5 >> 8) ))) << 8) | + (((unsigned int)(vmkKey[20] ^ ((unsigned char) (schedule5)))) << 0); + + schedule6= + (((unsigned int)(vmkKey[27] ^ ((unsigned char) (schedule6 >> 24) ))) << 24) | + (((unsigned int)(vmkKey[26] ^ ((unsigned char) (schedule6 >> 16) ))) << 16) | + (((unsigned int)(vmkKey[25] ^ ((unsigned char) (schedule6 >> 8) ))) << 8) | + (((unsigned int)(vmkKey[24] ^ ((unsigned char) (schedule6)))) << 0); + + schedule7= + (((unsigned int)(vmkKey[31] ^ ((unsigned char) (schedule7 >> 24) ))) << 24) | + (((unsigned int)(vmkKey[30] ^ ((unsigned char) (schedule7 >> 16) ))) << 16) | + (((unsigned int)(vmkKey[29] ^ ((unsigned char) (schedule7 >> 8) ))) << 8) | + (((unsigned int)(vmkKey[28] ^ ((unsigned char) (schedule7)))) << 0); + + + d += 0x01000000; + + encrypt( + hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7, + a, b, c, d, + &(schedule8), &(schedule9), &(schedule10), &(schedule11) + ); + + schedule8= + (((unsigned int)(vmkKey[35] ^ ((unsigned char) (schedule8 >> 24) ))) << 24) | + (((unsigned int)(vmkKey[34] ^ ((unsigned char) (schedule8 >> 16) ))) << 16) | + (((unsigned int)(vmkKey[33] ^ ((unsigned char) (schedule8 >> 8) ))) << 8) | + (((unsigned int)(vmkKey[32] ^ ((unsigned char) (schedule8)))) << 0); + + schedule9= + (((unsigned int)(vmkKey[39] ^ ((unsigned char) (schedule9 >> 24) ))) << 24) | + (((unsigned int)(vmkKey[38] ^ ((unsigned char) (schedule9 >> 16) ))) << 16) | + (((unsigned int)(vmkKey[37] ^ ((unsigned char) (schedule9 >> 8) ))) << 8) | + (((unsigned int)(vmkKey[36] ^ ((unsigned char) (schedule9)))) << 0); + + schedule10= + (((unsigned int)(vmkKey[43] ^ ((unsigned char) (schedule10 >> 24) ))) << 24) | + (((unsigned int)(vmkKey[42] ^ ((unsigned char) (schedule10 >> 16) ))) << 16) | + (((unsigned int)(vmkKey[41] ^ ((unsigned char) (schedule10 >> 8) ))) << 8) | + (((unsigned int)(vmkKey[40] ^ ((unsigned char) (schedule10)))) << 0); + + encrypt( + hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7, + macIV0, macIV4, macIV8, macIV12, + &(schedule16), &(schedule17), &(schedule18), &(schedule19) + ); + + encrypt( + hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7, + computeMacIV0, computeMacIV4, computeMacIV8, computeMacIV12, + &(schedule12), &(schedule13), &(schedule14), &(schedule15) + ); + + schedule28 = schedule0 ^ schedule12; + schedule29 = schedule1 ^ schedule13; + schedule30 = schedule2 ^ schedule14; + schedule31 = schedule3 ^ schedule15; + + encrypt( + hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7, + schedule28, schedule29, schedule30, schedule31, + &(schedule12), &(schedule13), &(schedule14), &(schedule15) + ); + + schedule28 = schedule4 ^ schedule12; + schedule29 = schedule5 ^ schedule13; + schedule30 = schedule6 ^ schedule14; + schedule31 = schedule7 ^ schedule15; + + encrypt( + hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7, + schedule28, schedule29, schedule30, schedule31, + &(schedule12), &(schedule13), &(schedule14), &(schedule15) + ); + + schedule28 = schedule8 ^ schedule12; + schedule29 = schedule9 ^ schedule13; + schedule30 = schedule10 ^ schedule14; + schedule31 = schedule15; + + encrypt( + hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7, + schedule28, schedule29, schedule30, schedule31, + &(schedule12), &(schedule13), &(schedule14), &(schedule15) + ); + + if ( + + ( + schedule12 == ( (unsigned int) + (((unsigned int)(mac[3] ^ ((unsigned char) (schedule16 >> 24) ))) << 24) | + (((unsigned int)(mac[2] ^ ((unsigned char) (schedule16 >> 16) ))) << 16) | + (((unsigned int)(mac[1] ^ ((unsigned char) (schedule16 >> 8) ))) << 8) | + (((unsigned int)(mac[0] ^ ((unsigned char) (schedule16)))) << 0) ) + ) + && + ( + schedule13 == ( (unsigned int) + (((unsigned int)(mac[7] ^ ((unsigned char) (schedule17 >> 24) ))) << 24) | + (((unsigned int)(mac[6] ^ ((unsigned char) (schedule17 >> 16) ))) << 16) | + (((unsigned int)(mac[5] ^ ((unsigned char) (schedule17 >> 8) ))) << 8) | + (((unsigned int)(mac[4] ^ ((unsigned char) (schedule17)))) << 0) ) + ) + && + ( + schedule14 == ( (unsigned int) + (((unsigned int)(mac[11] ^ ((unsigned char) (schedule18 >> 24) ))) << 24) | + (((unsigned int)(mac[10] ^ ((unsigned char) (schedule18 >> 16) ))) << 16) | + (((unsigned int)(mac[9] ^ ((unsigned char) (schedule18 >> 8) ))) << 8) | + (((unsigned int)(mac[8] ^ ((unsigned char) (schedule18)))) << 0) ) + ) + && + ( + schedule15 == ( (unsigned int) + (((unsigned int)(mac[15] ^ ((unsigned char) (schedule19 >> 24) ))) << 24) | + (((unsigned int)(mac[14] ^ ((unsigned char) (schedule19 >> 16) ))) << 16) | + (((unsigned int)(mac[13] ^ ((unsigned char) (schedule19 >> 8) ))) << 8) | + (((unsigned int)(mac[12] ^ ((unsigned char) (schedule19)))) << 0) ) + ) + ) + { + found[0] = gIndex; + break; + } + + gIndex += get_global_size(0); + } + + return; +} diff --git a/OpenCL_version/main.c b/OpenCL_version/main.c index fe1da93..99bbb12 100755 --- a/OpenCL_version/main.c +++ b/OpenCL_version/main.c @@ -40,7 +40,7 @@ int psw_x_thread=8; int tot_psw=0; size_t size_psw=0; size_t tot_word_mem=(SINGLE_BLOCK_SHA_SIZE * ITERATION_NUMBER * sizeof(uint32_t)); -int strict_check=0; +int strict_check=0, mac_comparison=0; // OpenCL Vars cl_context cxGPUContext; // OpenCL context @@ -62,10 +62,12 @@ void usage(char *name) "\t\tPath to dictionary or alphabet file\n" " -s" "\t\tStrict check (use only in case of false positives, faster solution)\n" + " -m" + "\t\tMAC comparison (use only in case of false positives, slower solution)\n" " -p" "\t\tPlatform\n" " -g" - "\t\tGPU device number\n" + "\t\tDevice number\n" " -t" "\t\tSet the number of password per thread threads\n" " -b" @@ -75,97 +77,97 @@ void usage(char *name) int checkDeviceStatistics() { int i, j; - char* value; - size_t valueSize, maxWorkGroup; - cl_uint platformCount; - cl_platform_id* platforms; - cl_uint deviceCount; - cl_device_id* devices; - cl_uint maxComputeUnits, deviceAddressBits; - cl_ulong maxAllocSize, maxConstBufSize; - cl_uint ccMajor, ccMinor, registersPerBlock, warpSize, overlap; - char dname[2048]; - int deviceFound=0; + char* value; + size_t valueSize, maxWorkGroup; + cl_uint platformCount; + cl_platform_id* platforms; + cl_uint deviceCount; + cl_device_id* devices; + cl_uint maxComputeUnits, deviceAddressBits; + cl_ulong maxAllocSize, maxConstBufSize; + cl_uint ccMajor, ccMinor, registersPerBlock, warpSize, overlap; + char dname[2048]; + int deviceFound=0; size_t avail, total; - // get all platforms - clGetPlatformIDs(0, NULL, &platformCount); - platforms = (cl_platform_id*) malloc(sizeof(cl_platform_id) * platformCount); - clGetPlatformIDs(platformCount, platforms, NULL); - - - for (i = 0; i < platformCount; i++) - { - // get all devices - clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, 0, NULL, &deviceCount); - devices = (cl_device_id*) malloc(sizeof(cl_device_id) * deviceCount); - clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, deviceCount, devices, NULL); - - printf("\n# Platform: %d, # Devices: %d\n", i, deviceCount); - // for each device print critical attributes - for (j = 0; j < deviceCount; j++) - { + // get all platforms + clGetPlatformIDs(0, NULL, &platformCount); + platforms = (cl_platform_id*) malloc(sizeof(cl_platform_id) * platformCount); + clGetPlatformIDs(platformCount, platforms, NULL); - // print device name - clGetDeviceInfo(devices[j], CL_DEVICE_NAME, 0, NULL, &valueSize); - value = (char*) malloc(valueSize); - clGetDeviceInfo(devices[j], CL_DEVICE_NAME, valueSize, value, NULL); - if (platform_id == i && gpu_id == j) - { - printf("\n====================================\nSelected device: %s (ID: %d) properties\n====================================\n\n", value, j); - deviceFound=1; - } - else - printf("\n====================================\nDevice %s (ID: %d) properties\n====================================\n\n", value, j); + for (i = 0; i < platformCount; i++) + { + // get all devices + clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, 0, NULL, &deviceCount); + devices = (cl_device_id*) malloc(sizeof(cl_device_id) * deviceCount); + clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, deviceCount, devices, NULL); - free(value); - - // print hardware device version - clGetDeviceInfo(devices[j], CL_DEVICE_VERSION, 0, NULL, &valueSize); - value = (char*) malloc(valueSize); - clGetDeviceInfo(devices[j], CL_DEVICE_VERSION, valueSize, value, NULL); - printf("OpenCL version supported: %s\n", value); - free(value); - - // print software driver version - clGetDeviceInfo(devices[j], CL_DRIVER_VERSION, 0, NULL, &valueSize); - value = (char*) malloc(valueSize); - clGetDeviceInfo(devices[j], CL_DRIVER_VERSION, valueSize, value, NULL); - printf("Software version: %s\n", value); - free(value); - - // print c version supported by compiler for device - clGetDeviceInfo(devices[j], CL_DEVICE_OPENCL_C_VERSION, 0, NULL, &valueSize); - value = (char*) malloc(valueSize); - clGetDeviceInfo(devices[j], CL_DEVICE_OPENCL_C_VERSION, valueSize, value, NULL); - printf("OpenCL C version: %s\n", value); - free(value); - - clGetDeviceInfo(devices[j], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(maxAllocSize), &maxAllocSize, NULL); - printf("Max Global Memory Size: %lld\n", maxAllocSize); - GPU_MAX_GLOBAL_MEM=maxAllocSize; - maxAllocSize=0; - clGetDeviceInfo(devices[j], CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(maxAllocSize), &maxAllocSize, NULL); - printf("Max Global Memory Alloc Size: %lld\n", maxAllocSize); - GPU_MAX_MEM_ALLOC_SIZE=maxAllocSize; + printf("\n# Platform: %d, # Devices: %d\n", i, deviceCount); + // for each device print critical attributes + for (j = 0; j < deviceCount; j++) + { - clGetDeviceInfo(devices[j], CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof(maxConstBufSize), &maxConstBufSize, NULL); - printf("Max Const Memory Buffer Size: %lld\n", maxConstBufSize); - clGetDeviceInfo(devices[j], CL_DEVICE_ADDRESS_BITS, sizeof(deviceAddressBits), &deviceAddressBits, NULL); - printf("Device Address Bits: %d\n", deviceAddressBits); + // print device name + clGetDeviceInfo(devices[j], CL_DEVICE_NAME, 0, NULL, &valueSize); + value = (char*) malloc(valueSize); + clGetDeviceInfo(devices[j], CL_DEVICE_NAME, valueSize, value, NULL); - // print parallel compute units - clGetDeviceInfo(devices[j], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(maxComputeUnits), &maxComputeUnits, NULL); - printf("Parallel compute units: %d\n", maxComputeUnits); - GPU_MAX_COMPUTE_UNITS=maxComputeUnits; + if (platform_id == i && gpu_id == j) + { + printf("\n====================================\nSelected device: %s (ID: %d) properties\n====================================\n\n", value, j); + deviceFound=1; + } + else + printf("\n====================================\nDevice %s (ID: %d) properties\n====================================\n\n", value, j); - clGetDeviceInfo(devices[j], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(maxWorkGroup), &maxWorkGroup, NULL); - printf("Max Workgroup Size: %zd\n", maxWorkGroup); - GPU_MAX_WORKGROUP_SIZE=maxWorkGroup; + free(value); + + // print hardware device version + clGetDeviceInfo(devices[j], CL_DEVICE_VERSION, 0, NULL, &valueSize); + value = (char*) malloc(valueSize); + clGetDeviceInfo(devices[j], CL_DEVICE_VERSION, valueSize, value, NULL); + printf("OpenCL version supported: %s\n", value); + free(value); + + // print software driver version + clGetDeviceInfo(devices[j], CL_DRIVER_VERSION, 0, NULL, &valueSize); + value = (char*) malloc(valueSize); + clGetDeviceInfo(devices[j], CL_DRIVER_VERSION, valueSize, value, NULL); + printf("Software version: %s\n", value); + free(value); + + // print c version supported by compiler for device + clGetDeviceInfo(devices[j], CL_DEVICE_OPENCL_C_VERSION, 0, NULL, &valueSize); + value = (char*) malloc(valueSize); + clGetDeviceInfo(devices[j], CL_DEVICE_OPENCL_C_VERSION, valueSize, value, NULL); + printf("OpenCL C version: %s\n", value); + free(value); - clGetDeviceInfo(devices[j], CL_DEVICE_VENDOR, sizeof(dname), dname, NULL); - printf("Vendor: %s\n", dname); + clGetDeviceInfo(devices[j], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(maxAllocSize), &maxAllocSize, NULL); + printf("Max Global Memory Size: %lld\n", maxAllocSize); + GPU_MAX_GLOBAL_MEM=maxAllocSize; + maxAllocSize=0; + clGetDeviceInfo(devices[j], CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(maxAllocSize), &maxAllocSize, NULL); + printf("Max Global Memory Alloc Size: %lld\n", maxAllocSize); + GPU_MAX_MEM_ALLOC_SIZE=maxAllocSize; + + clGetDeviceInfo(devices[j], CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof(maxConstBufSize), &maxConstBufSize, NULL); + printf("Max Const Memory Buffer Size: %lld\n", maxConstBufSize); + clGetDeviceInfo(devices[j], CL_DEVICE_ADDRESS_BITS, sizeof(deviceAddressBits), &deviceAddressBits, NULL); + printf("Device Address Bits: %d\n", deviceAddressBits); + + // print parallel compute units + clGetDeviceInfo(devices[j], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(maxComputeUnits), &maxComputeUnits, NULL); + printf("Parallel compute units: %d\n", maxComputeUnits); + GPU_MAX_COMPUTE_UNITS=maxComputeUnits; + + clGetDeviceInfo(devices[j], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(maxWorkGroup), &maxWorkGroup, NULL); + printf("Max Workgroup Size: %zd\n", maxWorkGroup); + GPU_MAX_WORKGROUP_SIZE=maxWorkGroup; + + clGetDeviceInfo(devices[j], CL_DEVICE_VENDOR, sizeof(dname), dname, NULL); + printf("Vendor: %s\n", dname); if(strstr(dname, "NVIDIA") != NULL) { @@ -297,7 +299,7 @@ int destroyClCtx() int main (int argc, char **argv) { char * input_dictionary=NULL, * input_hash=NULL; - unsigned char *salt, *nonce, *vmk; + unsigned char *salt, *nonce, *vmk, *mac; uint32_t * w_blocks_d; long int totGlobalMem; @@ -317,7 +319,7 @@ int main (int argc, char **argv) //*********************** Options ************************ while (1) { - opt = getopt(argc, argv, "hf:d:t:b:p:g:s"); + opt = getopt(argc, argv, "hf:d:t:b:p:g:ms"); if (opt == -1) break; switch (opt) { @@ -357,6 +359,9 @@ int main (int argc, char **argv) case 'g': gpu_id = atoi(optarg); break; + case 'm': + mac_comparison = 1; + break; case 'p': platform_id = atoi(optarg); break; @@ -403,11 +408,17 @@ int main (int argc, char **argv) //****************** Data from target file ******************* printf("\n====================================\nExtracting data from disk image\n====================================\n\n"); - if(parse_data(input_hash, &salt, &nonce, &vmk) == BIT_FAILURE) + if(parse_data(input_hash, &salt, &nonce, &vmk, &mac) == BIT_FAILURE) { fprintf(stderr, "Input hash format error... exit!\n"); goto cleanup; } + + if(mac_comparison == 1 && mac == NULL) + { + fprintf(stderr, "MAC comparison option selected but no MAC string found in input hash. MAC comparison not used!\n"); + mac_comparison=0; + } //************************************************************ printf("\n\n====================================\nDictionary attack\n====================================\n\n"); @@ -421,7 +432,7 @@ int main (int argc, char **argv) //********************************************** //************* Dictionary Attack ************* - opencl_attack(input_dictionary, w_blocks_h, vmk, nonce, gridBlocks); + opencl_attack(input_dictionary, w_blocks_h, vmk, nonce, mac, gridBlocks); //********************************************* cleanup: diff --git a/OpenCL_version/opencl_attack.c b/OpenCL_version/opencl_attack.c index bc35a75..af9f678 100755 --- a/OpenCL_version/opencl_attack.c +++ b/OpenCL_version/opencl_attack.c @@ -45,26 +45,34 @@ static int check_match() { } -char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryptedVMK, unsigned char * nonce, int gridBlocks) +char *opencl_attack(char *dname, unsigned int * w_blocks, + unsigned char * encryptedVMK, + unsigned char * nonce, unsigned char * encryptedMAC, + int gridBlocks) { cl_device_id device_id; cl_program cpProgram; - cl_kernel ckKernelAttack; - cl_mem deviceEncryptedVMK, deviceIV, devicePassword, deviceFound, w_blocks_d; + cl_kernel ckKernelAttack; + char vmkIV[IV_SIZE], macIV[IV_SIZE], computeMacIV[IV_SIZE]; + cl_mem d_vmk, d_mac, d_macIV, d_computeMacIV; + cl_mem devicePassword, deviceFound, w_blocks_d; cl_int ciErr1, ciErr2, ccMajor; size_t szGlobalWorkSize; size_t szLocalWorkSize; int numReadPassword, numPassword, passwordBufferSize, ret, totPsw=0; - char tmpIV[IV_SIZE]; FILE *fp_kernel, *fp_file_passwords; - //Very very ugly... + //Really ugly... char fileNameAttack[] = "./OpenCL_version/kernel_attack.cl"; size_t source_size_attack, source_size; char *source_str_attack; size_t len = 0; cl_int ret_cl = CL_SUCCESS, ret_cl_log = CL_SUCCESS, ret_info_kernel = CL_SUCCESS; char optProgram[128]; + + unsigned int vmkIV0, vmkIV4, vmkIV8, vmkIV12; + unsigned int macIV0, macIV4, macIV8, macIV12; + unsigned int cMacIV0, cMacIV4, cMacIV8, cMacIV12; //------- READ CL FILE ------ fp_kernel = fopen(fileNameAttack, "rb"); @@ -90,19 +98,51 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp fprintf(stderr, "crack_dict input error\n"); return NULL; } - - //-------- IV setup ------ - memset(tmpIV, 0, IV_SIZE); - memcpy(tmpIV + 1, nonce, NONCE_SIZE); - if(IV_SIZE-1 - NONCE_SIZE - 1 < 0) +/* + if(tot_psw <= 0) { - fprintf(stderr, "Nonce error\n"); + fprintf(stderr, "Attack tot passwords error: %d\n", tot_psw); return NULL; } - *tmpIV = (unsigned char)(IV_SIZE - 1 - NONCE_SIZE - 1); - tmpIV[IV_SIZE-1] = 1; - // -------------------------------------------------------------------------- +*/ + //-------- vmkIV setup ------ + memset(vmkIV, 0, IV_SIZE); + vmkIV[0] = (unsigned char)(IV_SIZE - 1 - NONCE_SIZE - 1); + memcpy(vmkIV + 1, nonce, NONCE_SIZE); + if(IV_SIZE-1 - NONCE_SIZE - 1 < 0) + { + fprintf(stderr, "Attack nonce error\n"); + return NULL; + } + vmkIV[IV_SIZE-1] = 1; + // ----------------------- + if(mac_comparison == 1) + { + //-------- macIV setup ------ + memset(macIV, 0, IV_SIZE); + macIV[0] = (unsigned char)(IV_SIZE - 1 - NONCE_SIZE - 1); + memcpy(macIV + 1, nonce, NONCE_SIZE); + if(IV_SIZE-1 - NONCE_SIZE - 1 < 0) + { + fprintf(stderr, "Attack nonce error\n"); + return NULL; + } + macIV[IV_SIZE-1] = 0; + // ----------------------- + + //-------- computeMacIV setup ------ + memset(computeMacIV, 0, IV_SIZE); + computeMacIV[0] = 0x3a; + memcpy(computeMacIV + 1, nonce, NONCE_SIZE); + if(IV_SIZE-1 - NONCE_SIZE - 1 < 0) + { + fprintf(stderr, "Attack nonce error\n"); + return NULL; + } + computeMacIV[IV_SIZE-1] = 0x2c; + // ----------------------- + } // ---- Open File Dictionary ---- if (!memcmp(dname, "-\0", 2)) { fp_file_passwords= stdin; @@ -141,8 +181,16 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp CL_ERROR(ciErr1); } - ckKernelAttack = clCreateKernel(cpProgram, "opencl_bitcracker_attack", &ciErr1); - CL_ERROR(ciErr1); + if(mac_comparison == 1) + { + ckKernelAttack = clCreateKernel(cpProgram, "opencl_bitcracker_attack_mac", &ciErr1); + CL_ERROR(ciErr1); + } + else + { + ckKernelAttack = clCreateKernel(cpProgram, "opencl_bitcracker_attack", &ciErr1); + CL_ERROR(ciErr1); + } size_t workgroup_size; ret_info_kernel = clGetKernelWorkGroupInfo(ckKernelAttack, cdDevices[gpu_id], CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &workgroup_size, NULL); @@ -173,7 +221,7 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp // -------------------------------------------------------------------------- // ------------------------------- Data setup ------------------------------- - deviceEncryptedVMK = clCreateBuffer(cxGPUContext, CL_MEM_READ_WRITE, VMK_DECRYPT_SIZE*sizeof(unsigned char), NULL, &ciErr1); + d_vmk = clCreateBuffer(cxGPUContext, CL_MEM_READ_WRITE, VMK_FULL_SIZE*sizeof(unsigned char), NULL, &ciErr1); CL_ERROR(ciErr1); devicePassword = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, passwordBufferSize*sizeof(unsigned char), NULL, &ciErr1); @@ -184,33 +232,64 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp w_blocks_d = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, SINGLE_BLOCK_SHA_SIZE * ITERATION_NUMBER * sizeof(unsigned int), NULL, &ciErr1); CL_ERROR(ciErr1); + + if(mac_comparison == 1) + { + d_mac = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, MAC_SIZE * sizeof(char), NULL, &ciErr1); + CL_ERROR(ciErr1); + + d_macIV = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, IV_SIZE * sizeof(char), NULL, &ciErr1); + CL_ERROR(ciErr1); + + d_computeMacIV = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, IV_SIZE * sizeof(char), NULL, &ciErr1); + CL_ERROR(ciErr1); + } + // -------------------------------------------------------------------------- // ------------------------------- Write buffers ------------------------------- - unsigned int tmp_global = ((unsigned int *)(tmpIV))[0]; - unsigned int IV0=(unsigned int )(((unsigned int )(tmp_global & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(tmp_global & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(tmp_global & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(tmp_global & 0x000000ff) << 24); - - tmp_global = ((unsigned int *)(tmpIV+4))[0]; - unsigned int IV4=(unsigned int )(((unsigned int )(tmp_global & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(tmp_global & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(tmp_global & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(tmp_global & 0x000000ff) << 24); - - tmp_global = ((unsigned int *)(tmpIV+8))[0]; - unsigned int IV8=(unsigned int )(((unsigned int )(tmp_global & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(tmp_global & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(tmp_global & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(tmp_global & 0x000000ff) << 24); - - tmp_global = ((unsigned int *)(tmpIV+12))[0]; - unsigned int IV12=(unsigned int )(((unsigned int )(tmp_global & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(tmp_global & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(tmp_global & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(tmp_global & 0x000000ff) << 24); + vmkIV0 = ((unsigned int *)(vmkIV))[0]; + vmkIV4 = ((unsigned int *)(vmkIV+4))[0]; + vmkIV8 = ((unsigned int *)(vmkIV+8))[0]; + vmkIV12 = ((unsigned int *)(vmkIV+12))[0]; + + if(mac_comparison == 1) + { + macIV0 = ((unsigned int *)(macIV))[0]; + macIV4 = ((unsigned int *)(macIV+4))[0]; + macIV8 = ((unsigned int *)(macIV+8))[0]; + macIV12 = ((unsigned int *)(macIV+12))[0]; + + cMacIV0 = ((unsigned int *)(computeMacIV))[0]; + cMacIV4 = ((unsigned int *)(computeMacIV+4))[0]; + cMacIV8 = ((unsigned int *)(computeMacIV+8))[0]; + cMacIV12 = ((unsigned int *)(computeMacIV+12))[0]; + } ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, w_blocks_d, CL_TRUE, 0, SINGLE_BLOCK_SHA_SIZE * ITERATION_NUMBER * sizeof(int), w_blocks, 0, NULL, NULL); CL_ERROR(ciErr1); - ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, deviceEncryptedVMK, CL_TRUE, 0, VMK_DECRYPT_SIZE*sizeof(char), encryptedVMK+16, 0, NULL, NULL); + ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, d_vmk, CL_TRUE, 0, VMK_FULL_SIZE*sizeof(char), encryptedVMK, 0, NULL, NULL); CL_ERROR(ciErr1); + + if(mac_comparison == 1) + { + ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, d_mac, CL_TRUE, 0, MAC_SIZE*sizeof(char), encryptedMAC, 0, NULL, NULL); + CL_ERROR(ciErr1); + + ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, d_macIV, CL_TRUE, 0, IV_SIZE*sizeof(char), macIV, 0, NULL, NULL); + CL_ERROR(ciErr1); + + ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, d_computeMacIV, CL_TRUE, 0, IV_SIZE*sizeof(char), computeMacIV, 0, NULL, NULL); + CL_ERROR(ciErr1); + } // ---------------------------------------------------------------------------- szLocalWorkSize = GPU_MAX_WORKGROUP_SIZE; szGlobalWorkSize = gridBlocks*szLocalWorkSize; //TOT THREADS - printf("Starting OpenCL attack:\n\tLocal Work Size: %d\n\tWork Group Number: %d\n\tGlobal Work Size: %zd\n\tPassword per thread: %d\n\tPassword per kernel: %d\n\tDictionary: %s\n\n", - szLocalWorkSize, gridBlocks, szGlobalWorkSize, psw_x_thread, numPassword, (fp_file_passwords == stdin)?"standard input":dname); + printf("Starting OpenCL attack:\n\tLocal Work Size: %zd\n\tWork Group Number: %d\n\tGlobal Work Size: %zd\n\tPassword per thread: %d\n\tPassword per kernel: %d\n\tDictionary: %s\n\tStrict Check (-s): %s\n\tMAC Comparison (-m): %s\n\t\n\n", + szLocalWorkSize, gridBlocks, szGlobalWorkSize, psw_x_thread, numPassword, (fp_file_passwords == stdin)?"standard input":dname, (strict_check == 1)?"Yes":"No", (mac_comparison == 1)?"Yes":"No"); int iter=0; while(!feof(fp_file_passwords)) @@ -233,23 +312,54 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp ciErr1 |= clSetKernelArg(ckKernelAttack, 2, sizeof(cl_mem), (void*)&deviceFound); CL_ERROR(ciErr1); - ciErr1 |= clSetKernelArg(ckKernelAttack, 3, sizeof(cl_mem), (void*)&deviceEncryptedVMK); + ciErr1 |= clSetKernelArg(ckKernelAttack, 3, sizeof(cl_mem), (void*)&d_vmk); CL_ERROR(ciErr1); ciErr1 |= clSetKernelArg(ckKernelAttack, 4, sizeof(cl_mem), (void*)&w_blocks_d); CL_ERROR(ciErr1); - ciErr1 |= clSetKernelArg(ckKernelAttack, 5, sizeof(cl_int), (void*)&IV0); + ciErr1 |= clSetKernelArg(ckKernelAttack, 5, sizeof(cl_int), (void*)&vmkIV0); CL_ERROR(ciErr1); - ciErr1 |= clSetKernelArg(ckKernelAttack, 6, sizeof(cl_int), (void*)&IV4); + ciErr1 |= clSetKernelArg(ckKernelAttack, 6, sizeof(cl_int), (void*)&vmkIV4); CL_ERROR(ciErr1); - ciErr1 |= clSetKernelArg(ckKernelAttack, 7, sizeof(cl_int), (void*)&IV8); + ciErr1 |= clSetKernelArg(ckKernelAttack, 7, sizeof(cl_int), (void*)&vmkIV8); CL_ERROR(ciErr1); - ciErr1 |= clSetKernelArg(ckKernelAttack, 8, sizeof(cl_int), (void*)&IV12); + ciErr1 |= clSetKernelArg(ckKernelAttack, 8, sizeof(cl_int), (void*)&vmkIV12); CL_ERROR(ciErr1); + + if(mac_comparison == 1) + { + ciErr1 |= clSetKernelArg(ckKernelAttack, 9, sizeof(cl_mem), (void*)&d_mac); + CL_ERROR(ciErr1); + + ciErr1 |= clSetKernelArg(ckKernelAttack, 10, sizeof(cl_int), (void*)&macIV0); + CL_ERROR(ciErr1); + + ciErr1 |= clSetKernelArg(ckKernelAttack, 11, sizeof(cl_int), (void*)&macIV4); + CL_ERROR(ciErr1); + + ciErr1 |= clSetKernelArg(ckKernelAttack, 12, sizeof(cl_int), (void*)&macIV8); + CL_ERROR(ciErr1); + + ciErr1 |= clSetKernelArg(ckKernelAttack, 13, sizeof(cl_int), (void*)&macIV12); + CL_ERROR(ciErr1); + + + ciErr1 |= clSetKernelArg(ckKernelAttack, 14, sizeof(cl_int), (void*)&cMacIV0); + CL_ERROR(ciErr1); + + ciErr1 |= clSetKernelArg(ckKernelAttack, 15, sizeof(cl_int), (void*)&cMacIV4); + CL_ERROR(ciErr1); + + ciErr1 |= clSetKernelArg(ckKernelAttack, 16, sizeof(cl_int), (void*)&cMacIV8); + CL_ERROR(ciErr1); + + ciErr1 |= clSetKernelArg(ckKernelAttack, 17, sizeof(cl_int), (void*)&cMacIV12); + CL_ERROR(ciErr1); + } // -------------------------------------------------------- time_t start,end; @@ -273,7 +383,6 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp printf("OpenCL Kernel execution #%d\n\tEffective number psw: %d\n\tTime: %f sec\n\tPasswords x second: %10.2f pw/sec\n", iter, numReadPassword, TIMER_ELAPSED(0)/1.0E+6, numReadPassword/(TIMER_ELAPSED(0)/1.0E+6)); - // printf ("TIME timer: %d passwords in %.2lf seconds => %.2f pwd/s\n", numReadPassword, dif, (double)(numReadPassword/dif) ); ret = clFlush(cqCommandQueue); ret = clFinish(cqCommandQueue); @@ -306,7 +415,7 @@ out: if(cpProgram)clReleaseProgram(cpProgram); if(w_blocks_d)clReleaseMemObject(w_blocks_d); if(devicePassword)clReleaseMemObject(devicePassword); - if(deviceEncryptedVMK)clReleaseMemObject(deviceEncryptedVMK); + if(d_vmk)clReleaseMemObject(d_vmk); if(deviceFound)clReleaseMemObject(deviceFound); free(source_str_attack); diff --git a/OpenCL_version/utils.c b/OpenCL_version/utils.c index 1565be2..7599e2e 100755 --- a/OpenCL_version/utils.c +++ b/OpenCL_version/utils.c @@ -73,7 +73,7 @@ void print_hex(unsigned char *str, int len) printf("%02x", str[i]); } -int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, unsigned char ** vmk) +int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, unsigned char ** vmk, unsigned char ** mac) { char * hash; char *p; @@ -81,10 +81,13 @@ int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, FILE * fphash; char tmp[2]; int j=0; + const char zero_string[17]="0000000000000000"; (*salt) = (unsigned char *) Calloc(SALT_SIZE, sizeof(unsigned char)); (*nonce) = (unsigned char *) Calloc(NONCE_SIZE, sizeof(unsigned char)); (*vmk) = (unsigned char *) Calloc(VMK_SIZE, sizeof(unsigned char)); + (*mac) = (unsigned char *) Calloc(MAC_SIZE, sizeof(unsigned char)); + hash = (char *) Calloc(INPUT_HASH_SIZE, sizeof(char)); if(!input_hash) @@ -164,7 +167,21 @@ int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, } p = strtokm(NULL, "$"); // data - for (i = 0, j = 0; i < vmk_size*2; i+=2, j++) + for (i = 0, j = 0; i < MAC_SIZE*2; i+=2, j++) + { + tmp[0] = p[i]; + tmp[1] = p[i+1]; + long int ret = strtol(tmp, NULL, 16); + (*mac)[j] = (unsigned char)(ret); //((ARCH_INDEX(p[i * 2]) * 16) + ARCH_INDEX(p[i * 2 + 1])); + } + + if(mac_comparison == 1 && !memcmp((*mac), zero_string, MAC_SIZE)) + { + free(*mac); + (*mac)=NULL; + } + + for (j=0; i < vmk_size*2; i+=2, j++) { tmp[0] = p[i]; tmp[1] = p[i+1]; @@ -178,6 +195,7 @@ int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, free(*salt); free(*nonce); free(*vmk); + free(*mac); return BIT_FAILURE; }