From e9c8f26e2d86b306f7d1c4a738fb74cfec433cfa Mon Sep 17 00:00:00 2001 From: elenago Date: Tue, 19 Sep 2017 10:59:03 +0200 Subject: [PATCH] Hash extractor (fully compatbile with JtR hash format), both CUDA and OpenCL versions changed: they no longer accept an encrypted memory unit, instead they require the hash extractor output. Strict check added (-s option) to avoid false positives (empirically verified, it works with this repo images encrypted with Windows 7, 8.1 and 10). --- CUDA_version/bitcracker.h | 14 +- CUDA_version/cuda_attack.cu | 93 ++----- CUDA_version/main.cu | 58 ++-- CUDA_version/utils.cu | 202 ++++++++------ Hash_Extractor/Makefile | 6 + Hash_Extractor/bitcracker_hash | Bin 0 -> 13808 bytes Hash_Extractor/bitcracker_hash.c | 290 ++++++++++++++++++++ Hash_Extractor/imgWin10Compatible.txt | 1 + OpenCL_version/bitcracker.h | 17 +- OpenCL_version/kernel_attack.cl | 371 ++------------------------ OpenCL_version/main.c | 50 ++-- OpenCL_version/opencl_attack.c | 48 ++-- OpenCL_version/utils.c | 203 ++++++++------ README.md | 68 +++-- build.sh | 17 ++ run.sh | 19 -- run_tests.sh | 40 +++ 17 files changed, 767 insertions(+), 730 deletions(-) create mode 100755 Hash_Extractor/Makefile create mode 100755 Hash_Extractor/bitcracker_hash create mode 100644 Hash_Extractor/bitcracker_hash.c create mode 100644 Hash_Extractor/imgWin10Compatible.txt delete mode 100755 run.sh create mode 100755 run_tests.sh diff --git a/CUDA_version/bitcracker.h b/CUDA_version/bitcracker.h index e4b026f..54d15fa 100755 --- a/CUDA_version/bitcracker.h +++ b/CUDA_version/bitcracker.h @@ -39,11 +39,14 @@ #define MAC_SIZE 16 #define NONCE_SIZE 12 #define IV_SIZE 16 -#define VMK_SIZE 44 +#define VMK_SIZE 60 #define VMK_DECRYPT_SIZE 16 #define DICT_BUFSIZE (50*1024*1024) #define MAX_PLEN 32 +#define HASH_TAG "$bitlocker$" +#define HASH_TAG_LEN (sizeof(HASH_TAG) - 1) +#define INPUT_HASH_SIZE 210 #ifndef UINT32_C #define UINT32_C(c) c ## UL #endif @@ -65,7 +68,6 @@ #define ATTACK_DEFAULT_THREADS 1024 - #define BIT_SUCCESS 0 #define BIT_FAILURE 1 @@ -88,17 +90,17 @@ extern int gpu_id; extern int psw_x_thread; extern int tot_psw; extern size_t size_psw; +extern int strict_check; /* ++++++++++++++++++++++++++++++++++++++ DEVICE FUNCTIONS ++++++++++++++++++++++++++++++++++++++ */ __global__ void w_block_evaluate(unsigned char salt[SALT_SIZE], int totNumIteration, unsigned char padding[PADDING_SIZE], uint32_t * w_blocks); -__global__ __launch_bounds__(1024,1) void decrypt_vmk(int numStream, int numPassword, int *found, unsigned char * vmkKey, unsigned char * IV); +__global__ __launch_bounds__(1024,1) void decrypt_vmk(int numStream, int numPassword, int *found, unsigned char * vmkKey, unsigned char * IV, int check); /* ++++++++++++++++++++++++++++++++++++++ HOST FUNCTIONS ++++++++++++++++++++++++++++++++++++++ */ int w_block_precomputed(unsigned char * salt, uint32_t * w_blocks_d); - int readFilePassword(char ** buf, int maxNumPsw, FILE *fp); -int readData(char * encryptedImagePath, unsigned char ** salt, unsigned char ** mac, unsigned char ** nonce, unsigned char ** encryptedVMK); - +int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, unsigned char ** vmk); +char *strtokm(char *s1, const char *delims); char *cuda_attack(char *dname, uint32_t * w_blocks_d, unsigned char * encryptedVMK, unsigned char * nonce, int gridBlocks); void setBufferPasswordSize(size_t avail, size_t * passwordBufferSize, int * numPassword); diff --git a/CUDA_version/cuda_attack.cu b/CUDA_version/cuda_attack.cu index f51fe73..bc85be4 100755 --- a/CUDA_version/cuda_attack.cu +++ b/CUDA_version/cuda_attack.cu @@ -102,7 +102,7 @@ char *cuda_attack(char *dname, uint32_t * w_blocks_d, unsigned char * encryptedV // ---- CUDA VARIABLES ---- BITCRACKER_CUDA_CHECK( cudaMalloc( (void ** ) &deviceEncryptedVMK, VMK_DECRYPT_SIZE*sizeof(uint8_t)) ); - BITCRACKER_CUDA_CHECK( cudaMemcpy(deviceEncryptedVMK, encryptedVMK, VMK_DECRYPT_SIZE*sizeof(uint8_t), cudaMemcpyHostToDevice) ); + BITCRACKER_CUDA_CHECK( cudaMemcpy(deviceEncryptedVMK, (encryptedVMK+16), VMK_DECRYPT_SIZE*sizeof(uint8_t), cudaMemcpyHostToDevice) ); BITCRACKER_CUDA_CHECK( cudaMalloc( (void ** ) &deviceIV, IV_SIZE*sizeof(uint8_t)) ); BITCRACKER_CUDA_CHECK( cudaMemcpy(deviceIV, tmpIV, IV_SIZE*sizeof(uint8_t), cudaMemcpyHostToDevice) ); @@ -148,7 +148,7 @@ char *cuda_attack(char *dname, uint32_t * w_blocks_d, unsigned char * encryptedV BITCRACKER_CUDA_CHECK( cudaStreamSynchronize(stream[indexStream^1]) ); BITCRACKER_CUDA_CHECK( cudaEventRecord(start[indexStream], stream[indexStream]) ); - decrypt_vmk<<>>(indexStream, numReadPassword[indexStream], deviceFound[indexStream], deviceEncryptedVMK, deviceIV); + decrypt_vmk<<>>(indexStream, numReadPassword[indexStream], deviceFound[indexStream], deviceEncryptedVMK, deviceIV, strict_check); BITCRACKER_CUDA_CHECK_LAST_ERROR(); BITCRACKER_CUDA_CHECK( cudaEventRecord(stop[indexStream], stream[indexStream]) ); BITCRACKER_CUDA_CHECK( cudaMemcpyAsync(hostFound[indexStream], deviceFound[indexStream], sizeof(unsigned int), cudaMemcpyDeviceToHost, stream[indexStream]) ); @@ -206,11 +206,9 @@ char *cuda_attack(char *dname, uint32_t * w_blocks_d, unsigned char * encryptedV #define END_STRING 0x80 //0xFF -//16 byte per password + 1 byte per length -__global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsigned char * vmkKey, unsigned char * IV) { - int globalIndexPassword = (threadIdx.x+blockIdx.x*blockDim.x); - - //Avoid register spilling in local memory +__global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsigned char * vmkKey, unsigned char * IV, int strict_check) { + + int gIndex = (threadIdx.x+blockIdx.x*blockDim.x); uint32_t hash0; uint32_t hash1; uint32_t hash2; @@ -264,11 +262,10 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig uint32_t first_hash6; uint32_t first_hash7; - uint32_t indexW=(globalIndexPassword*FIXED_PASSWORD_BUFFER); + uint32_t indexW=(gIndex*FIXED_PASSWORD_BUFFER); int8_t curr_fetch=0; - //int8_t stop=0; - while(globalIndexPassword < tot_psw_kernel) + while(gIndex < tot_psw_kernel) { first_hash0 = UINT32_C(0x6A09E667); @@ -290,10 +287,9 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig h = UINT32_C(0x5BE0CD19); //----------------------------------------------------- FIRST HASH ------------------------------------------------ - indexW=(globalIndexPassword*FIXED_PASSWORD_BUFFER); + indexW=(gIndex*FIXED_PASSWORD_BUFFER); curr_fetch=0; index_generic=MAX_INPUT_PASSWORD_LEN; - //stop=0; if(numStream == 0) { schedule0 = ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch+1)) << 8) | 0; @@ -350,11 +346,9 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig if(tex1Dfetch(w_password0, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; /* stop=1; */ } curr_fetch+=2; - //27 schedule13 = ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch+1)) << 8) | 0; if(tex1Dfetch(w_password0, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; /* stop=1; */ } if(tex1Dfetch(w_password0, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; /* stop=1; */ } - //curr_fetch+=2; } else { @@ -412,39 +406,18 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig if(tex1Dfetch(w_password1, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; /* stop=1; */ } curr_fetch+=2; - //27 schedule13 = ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch+1)) << 8) | 0; if(tex1Dfetch(w_password1, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; /* stop=1; */ } if(tex1Dfetch(w_password1, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; /* stop=1; */ } - //curr_fetch+=2; } if(index_generic == MAX_INPUT_PASSWORD_LEN) schedule13 = schedule13 | ((uint32_t)0x8000); - //64-bit + schedule14=0; index_generic*=2; schedule15 = ((uint8_t)((index_generic << 3) >> 8)) << 8 | ((uint8_t)(index_generic << 3)); -/* - printf("thread: %d, indexW: %d, index_generic %d, MAX_INPUT_PASSWORD_LEN %d\n", globalIndexPassword, indexW, index_generic, MAX_INPUT_PASSWORD_LEN); - printf("thread: %d schedule0: %x\n", globalIndexPassword, schedule0); - printf("thread: %d schedule1: %x\n", globalIndexPassword, schedule1); - printf("thread: %d schedule2: %x\n", globalIndexPassword, schedule2); - printf("thread: %d schedule3: %x\n", globalIndexPassword, schedule3); - printf("thread: %d schedule4: %x\n", globalIndexPassword, schedule4); - printf("thread: %d schedule5: %x\n", globalIndexPassword, schedule5); - printf("thread: %d schedule6: %x\n", globalIndexPassword, schedule6); - printf("thread: %d schedule7: %x\n", globalIndexPassword, schedule7); - printf("thread: %d schedule8: %x\n", globalIndexPassword, schedule8); - printf("thread: %d schedule9: %x\n", globalIndexPassword, schedule9); - printf("thread: %d schedule10: %x\n", globalIndexPassword, schedule10); - printf("thread: %d schedule11: %x\n", globalIndexPassword, schedule11); - printf("thread: %d schedule12: %x\n", globalIndexPassword, schedule12); - printf("thread: %d schedule13: %x\n", globalIndexPassword, schedule13); - printf("thread: %d schedule14: %x\n", globalIndexPassword, schedule14); - printf("thread: %d schedule15: %x\n", globalIndexPassword, schedule15); -*/ ALL_SCHEDULE_LAST16() ROUND(a, b, c, d, e, f, g, h, schedule0, 0x428A2F98) @@ -642,7 +615,7 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig first_hash6 += g; first_hash7 += h; -//----------------------------------------------------- LOOP HASH ------------------------------------------------ + //----------------------------------------------------- LOOP HASH ------------------------------------------------ hash0=0; hash1=0; @@ -657,7 +630,6 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig //#pragma unroll 1048576 for(index_generic=0; index_generic < ITERATION_NUMBER/2; index_generic++) { - //Prima parte a = UINT32_C(0x6A09E667); b = UINT32_C(0xBB67AE85); c = UINT32_C(0x3C6EF372); @@ -764,7 +736,6 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig hash6 = UINT32_C(0x1F83D9AB) + g; hash7 = UINT32_C(0x5BE0CD19) + h; - //Seconda parte a = hash0; b = hash1; c = hash2; @@ -774,7 +745,6 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig g = hash6; h = hash7; - //I primi 4 valori dei blocchi W sono sempre uguali 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) @@ -855,7 +825,6 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig for(index_generic=ITERATION_NUMBER/2; index_generic < ITERATION_NUMBER; index_generic++) { - //Prima parte a = UINT32_C(0x6A09E667); b = UINT32_C(0xBB67AE85); c = UINT32_C(0x3C6EF372); @@ -962,7 +931,6 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig hash6 = UINT32_C(0x1F83D9AB) + g; hash7 = UINT32_C(0x5BE0CD19) + h; - //Seconda parte a = hash0; b = hash1; c = hash2; @@ -972,7 +940,6 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig g = hash6; h = hash7; - //I primi 4 valori dei blocchi W sono sempre uguali 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) @@ -1050,13 +1017,7 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig indexW += SINGLE_BLOCK_W_SIZE; } -//----------------------------------------------------- EXP KEY 256 ------------------------------------------------ - // with Nb=4 and Nk=256 -> Nr=15 so (14+1)*Nb=60 32b-words are needed - // 48 words - //----- 1 - // AddRoundKey - /* REUSE OF SCHEDULE VARIABLES */ schedule0 = __byte_perm(((uint32_t *)(IV))[0], 0, 0x0123) ^ hash0; schedule1 = __byte_perm(((uint32_t *)(IV+4))[0], 0, 0x0123) ^ hash1; schedule2 = __byte_perm(((uint32_t *)(IV+8))[0], 0, 0x0123) ^ hash2; @@ -1091,7 +1052,6 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6); schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7); - //----- 2 hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^ (TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^ (TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^ @@ -1117,7 +1077,6 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7); - //----- 3 hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^ (TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^ (TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^ @@ -1143,8 +1102,6 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6); schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7); - - //----- 4 hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^ (TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^ (TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^ @@ -1169,8 +1126,6 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6); schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7); - - //----- 5 hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^ (TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^ (TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^ @@ -1196,7 +1151,6 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7); - //----- 6 hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^ (TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^ (TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^ @@ -1221,7 +1175,6 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6); schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7); - // last 4 words hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^ (TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^ (TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^ @@ -1230,7 +1183,6 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig hash2 ^= hash1; hash3 ^= hash2; - // NR-th round schedule0 = (TS2[(schedule4 >> 24) ] & 0xFF000000) ^ (TS3[(schedule5 >> 16) & 0xFF] & 0x00FF0000) ^ (TS0[(schedule6 >> 8) & 0xFF] & 0x0000FF00) ^ @@ -1256,20 +1208,31 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig schedule6 = __byte_perm(schedule2, 0, 0x0123); schedule7 = __byte_perm(schedule3, 0, 0x0123); -// ------ TEST VMK ------ if ( - ((vmkKey[0] ^ ((uint8_t) schedule4)) == VMK_SIZE) && + ((vmkKey[0] ^ ((uint8_t) schedule4)) == 0x2c) && ((vmkKey[1] ^ ((uint8_t) (schedule4 >> 8))) == 0x00) && - ((vmkKey[8] ^ ((uint8_t) schedule6)) <= 0x05) && ((vmkKey[9] ^ ((uint8_t) (schedule6 >> 8))) == 0x20) ) { - *found = globalIndexPassword; - break; + if( + (strict_check == 0 && ((vmkKey[8] ^ ((uint8_t) schedule6)) <= 0x05)) + || + (strict_check == 1 && ((vmkKey[8] ^ ((uint8_t) schedule6)) == 0x03)) + ) + { + printf("schedule4=%x, char: %x, >>8: %x\n", schedule4, (uint8_t)schedule4, (uint8_t) (schedule4 >> 8)); + printf("(vmkKey[0] ^ ((uint8_t) schedule4)): %x\n", (vmkKey[0] ^ ((uint8_t) schedule4))); + printf("(vmkKey[1] ^ ((uint8_t) (schedule4 >> 8))): %x\n", (vmkKey[1] ^ ((uint8_t) (schedule4 >> 8)))); + printf("schedule6=%x, char: %x, >>8: %x\n", schedule6, (uint8_t)schedule6, (uint8_t) (schedule6 >> 8)); + printf("(vmkKey[8] ^ ((uint8_t) schedule6)): %x\n", (vmkKey[8] ^ ((uint8_t) schedule6))); + printf("(vmkKey[9] ^ ((uint8_t) (schedule6 >> 8))): %x\n", (vmkKey[9] ^ ((uint8_t) (schedule6 >> 8)))); + + *found = gIndex; + break; + } } -// ------ LOOP ------ - globalIndexPassword += (blockDim.x * gridDim.x); + gIndex += (blockDim.x * gridDim.x); } return; diff --git a/CUDA_version/main.cu b/CUDA_version/main.cu index a12de1a..ffa5a0d 100755 --- a/CUDA_version/main.cu +++ b/CUDA_version/main.cu @@ -26,16 +26,19 @@ 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; void usage(char *name){ - printf("\nUsage: %s -i -d \n\n" + printf("\nUsage: %s -f -d \n\n" "Options:\n\n" " -h, --help" "\t\t\tShow this help\n" - " -i, --diskimage" - "\t\tPath to your disk image\n" + " -f, --hashfile" + "\t\tPath to your input hash file (HashExtractor output)\n" " -d, --dictionary file" "\t\tPath to dictionary or alphabet file\n" + " -s, --strict" + "\t\tStrict check (use only in case of false positives)\n" " -g, --gpu" "\t\tGPU device number\n" " -t, --passthread" @@ -97,14 +100,10 @@ int getGPUStats() int main (int argc, char **argv) { - char dictionaryFile[INPUT_SIZE]; - char diskImageFile[INPUT_SIZE]; - unsigned char *salt; - unsigned char *mac; - unsigned char *nonce; - unsigned char *encryptedVMK; + char * input_dictionary=NULL, * input_hash=NULL; + unsigned char *salt, *nonce, *vmk; uint32_t * w_blocks_d; - + int gridBlocks = 1; int opt, option_index = 0; @@ -121,7 +120,7 @@ int main (int argc, char **argv) static struct option long_options[] = { {"help", no_argument, 0, 'h'}, - {"diskimage", required_argument, 0, 'i'}, + {"hashfile", required_argument, 0, 'f'}, {"dictionary", required_argument, 0, 'd'}, //{"info", required_argument, 0, 'o'}, //{"cuda.threads", required_argument, 0, 't'}, @@ -131,21 +130,22 @@ int main (int argc, char **argv) {0, 0, 0, 0} }; - opt = getopt_long (argc, argv, "hi:d:t:b:g:", long_options, &option_index); + opt = getopt_long(argc, argv, "hf:d:t:b:g:s", long_options, &option_index); if (opt == -1) break; switch (opt) { case 'h': usage(argv[0]); - exit(1); + exit(EXIT_FAILURE); break; - case 'i': + case 'f': if(strlen(optarg) >= INPUT_SIZE) { - fprintf(stderr, "ERROR: Disk image path is bigger than %d\n", INPUT_SIZE); + fprintf(stderr, "ERROR: Inut hash file path is bigger than %d\n", INPUT_SIZE); exit(EXIT_FAILURE); } - strncpy(diskImageFile,optarg, strlen(optarg)+1); + input_hash=(char *)Calloc(INPUT_SIZE, sizeof(char)); + strncpy(input_hash, optarg, strlen(optarg)+1); break; case 'd': if(strlen(optarg) >= INPUT_SIZE) @@ -153,7 +153,8 @@ int main (int argc, char **argv) fprintf(stderr, "ERROR: Dictionary file path is bigger than %d\n", INPUT_SIZE); exit(EXIT_FAILURE); } - strncpy(dictionaryFile,optarg, strlen(optarg)+1); + input_dictionary=(char *)Calloc(INPUT_SIZE, sizeof(char)); + strncpy(input_dictionary,optarg, strlen(optarg)+1); break; case 't': psw_x_thread = atoi(optarg); @@ -169,6 +170,9 @@ int main (int argc, char **argv) case 'g': gpu_id = atoi(optarg); break; + case 's': + strict_check = 1; + break; default: exit(EXIT_FAILURE); } @@ -182,20 +186,19 @@ int main (int argc, char **argv) exit(EXIT_FAILURE); } - if (dictionaryFile == NULL){ + if (input_dictionary == NULL){ printf("Missing dictionary file!\n"); usage(argv[0]); exit(EXIT_FAILURE); } - if (diskImageFile[0] == '\0'){ - printf("Missing dick image file!\n"); + if (input_hash == NULL){ + printf("Missing input hash file!\n"); usage(argv[0]); exit(EXIT_FAILURE); } //*********************************************************** - tot_psw=(ATTACK_DEFAULT_THREADS*gridBlocks*psw_x_thread); size_psw = tot_psw * FIXED_PASSWORD_BUFFER * sizeof(uint8_t); //****************** GPU device ******************* @@ -208,14 +211,10 @@ int main (int argc, char **argv) //****************** Data from target file ******************* printf("\n====================================\nExtracting data from disk image\n====================================\n\n"); - salt = (unsigned char *) Calloc(SALT_SIZE, sizeof(unsigned char)); - mac = (unsigned char *) Calloc(MAC_SIZE, sizeof(unsigned char)); - nonce = (unsigned char *) Calloc(NONCE_SIZE, sizeof(unsigned char)); - encryptedVMK = (unsigned char *) Calloc(VMK_SIZE, sizeof(unsigned char)); - if(readData(diskImageFile, &salt, &mac, &nonce, &encryptedVMK) == BIT_FAILURE) + if(parse_data(input_hash, &salt, &nonce, &vmk) == BIT_FAILURE) { - fprintf(stderr, "Disk image error... exit!\n"); + fprintf(stderr, "Input hash format error... exit!\n"); goto cleanup; } //************************************************************ @@ -231,14 +230,11 @@ int main (int argc, char **argv) //********************************************** //************* Dictionary Attack ************* - cuda_attack(dictionaryFile, w_blocks_d, encryptedVMK, nonce, gridBlocks); + cuda_attack(input_dictionary, w_blocks_d, vmk, nonce, gridBlocks); //********************************************* cleanup: - BITCRACKER_CUDA_CHECK( cudaFree(w_blocks_d) ); - printf("\n"); - return 0; } diff --git a/CUDA_version/utils.cu b/CUDA_version/utils.cu index 0b2aed4..66d42dd 100755 --- a/CUDA_version/utils.cu +++ b/CUDA_version/utils.cu @@ -21,6 +21,25 @@ #include "bitcracker.h" +/* John The Ripper function */ +char *strtokm(char *s1, const char *delims) +{ + static char *last = NULL; + char *endp; + + if (!s1) + s1 = last; + if (!s1 || *s1 == 0) + return last = NULL; + endp = strpbrk(s1, delims); + if (endp) { + *endp = '\0'; + last = endp + 1; + } else + last = NULL; + return s1; +} + void * Calloc(size_t len, size_t size) { void * ptr = NULL; if( size <= 0) @@ -54,109 +73,116 @@ void print_hex(unsigned char *str, int len) printf("%02x", str[i]); } -int readData(char * encryptedImagePath, unsigned char ** salt, unsigned char ** mac, unsigned char ** nonce, unsigned char ** encryptedVMK) +int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, unsigned char ** vmk) { - int match = 0; - const char signature[9] = "-FVE-FS-"; - int version = 0; - unsigned char vmk_entry[4] = { 0x02, 0x00, 0x08, 0x00 }; - unsigned char key_protection_type[2] = { 0x00, 0x20 }; - unsigned char value_type[2] = { 0x00, 0x05 }; - char c; - int i = 0; - int j, fileLen; + char * hash; + char *p; + int i, salt_len, iterations, vmk_size, nonce_len; + FILE * fphash; + char tmp[2]; + int j=0; - if( !salt || !mac || !nonce || !encryptedVMK ) { - fprintf(stderr, "Input error\n"); - return BIT_FAILURE; + (*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)); + hash = (char *) Calloc(INPUT_HASH_SIZE, sizeof(char)); + + if(!input_hash) + { + fprintf(stderr, "No input hash provided\n"); + goto out; } - printf("Opening file %s\n", encryptedImagePath); - FILE * encryptedImage = fopen(encryptedImagePath, "r"); - if (!encryptedImage) { - fprintf(stderr, "! %s : %s\n", encryptedImagePath, strerror(errno)); - return BIT_FAILURE; + fphash = fopen(input_hash, "r"); + if (!fphash) { + fprintf(stderr, "! %s : %s\n", input_hash, strerror(errno)); + goto out; + } + fgets(hash, INPUT_HASH_SIZE, fphash); + fclose(fphash); + if(!hash) + { + fprintf(stderr, "No correct input hash provided\n"); + goto out; } - fseek(encryptedImage, 0, SEEK_END); - fileLen = ftell(encryptedImage); - fseek(encryptedImage, 0, SEEK_SET); - for (j = 0; j < fileLen; j++) { - c = fgetc(encryptedImage); - while (i < 8 && (unsigned char)c == signature[i]) { - c = fgetc(encryptedImage); - i++; - } - if (i == 8) { - match = 1; - fprintf(stderr, "Signature found at 0x%08lx\n", (ftell(encryptedImage) - i - 1)); - fseek(encryptedImage, 1, SEEK_CUR); - version = fgetc(encryptedImage); - fprintf(stderr, "Version: %d ", version); - if (version == 1) - fprintf(stderr, "(Windows Vista)\n"); - else if (version == 2) - fprintf(stderr, "(Windows 7 or later)\n"); - else { - fprintf(stderr, "\nInvalid version, looking for a signature with valid version...\n"); - } - } - i = 0; - while (i < 4 && (unsigned char)c == vmk_entry[i]) { - c = fgetc(encryptedImage); - i++; - } + printf("Hash file %s:\n%s\n", input_hash, hash); - if (i == 4) { - fprintf(stderr, "VMK entry found at 0x%08lx\n", (ftell(encryptedImage) - i - 3)); - fseek(encryptedImage, 27, SEEK_CUR); - if ( - ( (unsigned char)fgetc(encryptedImage) == key_protection_type[0]) && - ( (unsigned char)fgetc(encryptedImage) == key_protection_type[1]) - ) - { - fprintf(stderr, "Key protector with user password found\n"); - fseek(encryptedImage, 12, SEEK_CUR); - fillBuffer(encryptedImage, *salt, SALT_SIZE); - fseek(encryptedImage, 83, SEEK_CUR); - if (((unsigned char)fgetc(encryptedImage) != value_type[0]) || ((unsigned char)fgetc(encryptedImage) != value_type[1])) { - fprintf(stderr, "Error: VMK not encrypted with AES-CCM\n"); - //ret failure? - } - - fseek(encryptedImage, 3, SEEK_CUR); - fillBuffer(encryptedImage, *nonce, NONCE_SIZE); - fillBuffer(encryptedImage, *mac, MAC_SIZE); - fillBuffer(encryptedImage, *encryptedVMK, VMK_SIZE); - - fprintf(stdout, "Nonce:\n"); - print_hex(*nonce, NONCE_SIZE); - fprintf(stdout, "\n"); - - fprintf(stdout, "MAC:\n"); - print_hex(*mac, MAC_SIZE); - fprintf(stdout, "\n"); - - fprintf(stdout, "VMK:\n"); - print_hex(*encryptedVMK, VMK_SIZE); - fprintf(stdout, "\n"); - break; - } - } - - i = 0; + if (strncmp(hash, HASH_TAG, HASH_TAG_LEN) != 0) + { + fprintf(stderr, "Wrong hash format\n"); + goto out; } - fclose(encryptedImage); + hash += HASH_TAG_LEN; + + p = strtokm(hash, "$"); // version + p = strtokm(NULL, "$"); // salt length - if (match == 0) { - fprintf(stderr, "Error while extracting data: No signature found!\n"); - return BIT_FAILURE; + salt_len = atoi(p); + if(salt_len != SALT_SIZE) + { + fprintf(stderr, "Wrong Salt size\n"); + goto out; } + p = strtokm(NULL, "$"); // salt + for (i = 0, j = 0; i < salt_len*2; i+=2, j++) + { + tmp[0] = p[i]; + tmp[1] = p[i+1]; + long int ret = strtol(tmp, NULL, 16); + (*salt)[j] = (unsigned char)(ret); //((ARCH_INDEX(p[i * 2]) * 16) + ARCH_INDEX(p[i * 2 + 1])); + } + + p = strtokm(NULL, "$"); // iterations + iterations = atoi(p); + p = strtokm(NULL, "$"); // nonce length + nonce_len = atoi(p); + if(nonce_len != NONCE_SIZE) + { + fprintf(stderr, "Wrong Nonce size\n"); + goto out; + } + + p = strtokm(NULL, "$"); // nonce + for (i = 0, j = 0; i < NONCE_SIZE*2; i+=2, j++) + { + tmp[0] = p[i]; + tmp[1] = p[i+1]; + long int ret = strtol(tmp, NULL, 16); + (*nonce)[j] = (unsigned char)(ret); //((ARCH_INDEX(p[i * 2]) * 16) + ARCH_INDEX(p[i * 2 + 1])); + } + + p = strtokm(NULL, "$"); // data_size + + vmk_size = atoi(p); + if(vmk_size != VMK_SIZE) + { + fprintf(stderr, "Wrong VMK size\n"); + goto out; + } + + p = strtokm(NULL, "$"); // data + for (i = 0, j = 0; i < vmk_size*2; i+=2, j++) + { + tmp[0] = p[i]; + tmp[1] = p[i+1]; + long int ret = strtol(tmp, NULL, 16); + (*vmk)[j] = (unsigned char)(ret); //((ARCH_INDEX(p[i * 2]) * 16) + ARCH_INDEX(p[i * 2 + 1])); + } + return BIT_SUCCESS; + + out: + free(*salt); + free(*nonce); + free(*vmk); + + return BIT_FAILURE; } + int readFilePassword(char ** buf, int maxNumPsw, FILE *fp) { int i=0, size; char tmp[FIXED_PASSWORD_BUFFER]; diff --git a/Hash_Extractor/Makefile b/Hash_Extractor/Makefile new file mode 100755 index 0000000..5472c3a --- /dev/null +++ b/Hash_Extractor/Makefile @@ -0,0 +1,6 @@ +bitcracker_hash: + gcc bitcracker_hash.c -o bitcracker_hash + +clean: + rm -rf *.o + rm -rf bitcracker_hash diff --git a/Hash_Extractor/bitcracker_hash b/Hash_Extractor/bitcracker_hash new file mode 100755 index 0000000000000000000000000000000000000000..c9bd163220311c87123ff8980c36364b8e9cba53 GIT binary patch literal 13808 zcmeHOeQ;FQb-x0X8U?%}#(-lCk5qDSie!-@TZoS!qfH;N00ja52;{NaT}f+JyPMs& zM!3eKSiDBh>rE5~+)U!O>P$YECgV}+Krs_xVN(m&nUr8phH2B5wxoUQ;FhGWgWbmT z_q+G*ry%2*{?Y!SH+uJ;bMHC7d+s^+-gn=9tG|Bt&G&vwh^m=F49*Zj_;6pY7Qzt* zC?RgaZQ{19&bHt4Kj`1Km8IoZRXKZ2jyTKFfMsp>-@m&`|}SgyggbX&*v3x(0E1r zh-dYMqls`5LX~*)Z&i3>8h||gUbv>%7fjj7NUT>2RN|f1cxMqnUgS%!GHg^>Rxln* z6&+ZKcW9}yZ;=KeUx`tQZ&~TszDO)&MPfa1iC-Sy?`ph!$jSdEJdTs|2%u%Pw{34L zX~|G_iCSF-US68#IH3zTR(M%fENZ3p_jkvmR>DqR7cURT@upoU^%>7Xoy6;n7ri#s zzFO#W+#!<}`AWPG={jbX6%G6*`_DEIYTuZl@GfW( z@@4Zvm*+Ff>P{7B#Zyf${ zX4m05l(8RTPK~@jmgiZ9YdQKW=VE?Bs$$tr3uqRzk2BW(1m;&2+eZ6c`%`wfzww?% z*b(VginN{WO~4<_@6bH`Gkr~aVA=!I9+>vP|6dR6Fx{V-LzfKGJ#fWz)0fRjzej7$ zp#ztN?+f)Lv(0Z{_iJXV@1D=BWlVU#!t?0`K1W^uMeiIKGXwoh`m*-|8Q1Rtz!N@4 z1H7sL-o7ahaI6S$!TYiVpw4My_e(vv8huL-(hNcS+IDuCpTBe&&6(NiZvbiKD<}A@ z>B#!ewmr|MmEZlVj_e;Txq9`gnYqoCjofOwyFJq#-@romXbZAQXBASSPC88mYDIxs zUZCnU)t;IA8g^OjZq`6ahx-rik4*RQ<@ZRR>u?YQHhT-gUQ1fNEYr$ChXN)4>F)3sQ4^zT-a#ii>Uu^Iln%povPT40MrbnuL z4mXfFY@Z5k{44f+`8F~RwTC9ETuJ1H$*DGd&RUM#O!YfxZPMY;qFm!MQP@3>+Pc~$6_l6eshegsm_OW*$Jgo2Xpbe(Si#_J}Aayj|Tl{3N2d?It0Qv#T`LuP2447%_Pw0zEQOD8q|zErDT12r}G{|h@y*ZC@B za=Rg-v}P3KcPaS~7UW-*@`ord-Is0Nszf}wC~!GB0ukjJ00Ml3!X#xfG%n|(&w)I| z=xJH^23}x_Y%$EUhO=K{53=lm@0i0uU5_bHI&sj~koOER6vFuV^ zs>59?hfj5foZO7f9H6<5jplL>Ve%H?xU@}$yR97VA{+?6{F+OKd~%FqA#j}p;TDv^ zab=tecWXJ^Nd@<80WK!tqJ(=9bD?M*7uKn8x0J(;DmbeEw^YLQ5iV8+$90)-U8+?3 zZlNw+hf+j$NWs+>;LhT8^LY&RDe(-U1@Aa&q4fYKz|Cq}3ZBcJM^i*row_0ft zmHts$y6#;#Li<p$Z=NF(IxZf5+K@lf*408GEkoc5~G z79Lo;WcORV7P9LE4DvnxF(|?_e+UasumocDN(_$qSrKHyn*ej1BJSjIkXS^$uly1r zm^=kP{E53tC_9kp;>$+P1cgRQjSeW}4bfJ$s+T2e`eU#KWQe3cdXGndD)Jal1NqwJG zh}55v`W&fIQjd{(npBq5lcc^z>Tyy>NHGk3KOhw&^(v_lslOrhAgP~|>IAhh8L=b5 zK-35)lksG$(I4)QC-)nv$RlARk}{fVM5V-LB@u{55T>KbTVx8#M7~m5wMM>f%>MF=s=Cw8BV5v0m+admfwf3LGgVl zV`n5~2kxjT7Tj&blSVXPhm$C+*%W&?5RHV4hn4aQBN~tIl?Fg60ORxR?TgrbMwx75 zVC} ztb5A-Hh)*c#*Le`n6a;qk;KQ+WFTnMvHU>PYHW#@_*S|Oz=(4)oJvP+qc4!^Yptp4 zj@TF)d&9}Prn>r2onVsBMUB`erbzIMr=T_mkO)Q4W;_;1?icKZ-IojpLZX|PDbbrq zi@r!}yjA$OZQHu7)z}nEq-`V8ALtFE(+D?4ZFi)%7e2820x^s$bgEP|p0?#Uy`FG= z)5ixy&5jfRIev|Xh_Tk69}M~Vk&Z>|^+rS7SPM8BK*;qqHCq!jJ=I!MV;Bv6D~yH) z_HFj;uD;rUpFZ#y>#179_nW9=7Qd6^ryc^SCz$m3PTZ{7SR#GBu#d6C9G1K*2g z;#(0lg&2kGC}fw38LB{*zQ3eSe{XKod5g~P(Rqi?cj|nP&L7gbt@FopKCJVnbpDLa zzoqjrouAP8CcFk+%V@bqqYGPfeWgr~?%sXZspm4ytNS|je5zfOI(suuQj>U(wl zbT;jQX%9?$VA=!I9+>vPv=fRaIY2Y+iUfv1Idv1bh36tm4|S|m=Jk}=O$$037m-ag5nQenPmT%DI$yhM4pQftRiQJqtbw3a=<{}tY{p+ zzlZn;%W%4y3d=^nO120#N)|uXFZAt*E{Mvum48K*h0~l_R(~IC4+qnB82>8}psViM zsCw3#2qaTsOP+>_PxJ>OG11Nwo|$57A(FF;kf@(P6LDv!>#79E%D{&(yDX&(frF=&P&rb%W z^$%6>=?eZ}1^;9P|5OG4wF>@ED|jB>m$v`C3jTNn|3U@-!wUZO3jU`R{KX3X-zxaG zEBLwK@%shAPnr1Loxm>yyn!qVg4y2jkFm8ca=7OBNck#gyhWO0SbuPtg- z$*(7RCDCh0VfE1KhF&nrR*L)*kzW_{YeaEzAjMT-%6e(x3_5yGx~B&Vf>OisFd(cn tkFrrxTpAh+Ur2=XP9qm7w#c(iYly literal 0 HcmV?d00001 diff --git a/Hash_Extractor/bitcracker_hash.c b/Hash_Extractor/bitcracker_hash.c new file mode 100644 index 0000000..b89b52a --- /dev/null +++ b/Hash_Extractor/bitcracker_hash.c @@ -0,0 +1,290 @@ +/* + * BitCracker: BitLocker password cracking tool, Hash Extractor. + * Copyright (C) 2013-2017 Elena Ago + * Massimo Bernaschi + * + * This file is part of BitCracker. + * + * BitCracker is free software: you can redistribute it and/or modify + * it under the terms of the GNU General Public License as published by + * the Free Software Foundation, either version 2 of the License, or + * (at your option) any later version. + * + * BitCracker is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with BitCracker. If not, see . + */ + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#define INPUT_SIZE 2048 + +#define SALT_SIZE 16 +#define MAC_SIZE 16 +#define NONCE_SIZE 12 +#define IV_SIZE 16 +#define VMK_SIZE 44 + +#define SIGNATURE_LEN 9 +#define VMK_SALT_JUMP 12 +#define VMK_AES_TYPE 0x0005 + +static unsigned char salt[SALT_SIZE], nonce[NONCE_SIZE], mac[MAC_SIZE], vmk[VMK_SIZE]; + +void * Calloc(size_t len, size_t size) { + void * ptr = NULL; + if( size <= 0) + { + fprintf(stderr,"Critical error: memory size is 0\n"); + exit(EXIT_FAILURE); + } + + ptr = (void *)calloc(len, size); + if( ptr == NULL ) + { + fprintf(stderr,"Critical error: Memory allocation\n"); + exit(EXIT_FAILURE); + } + return ptr; +} + + +static int usage(char *name){ + printf("\nUsage: %s -i -o \n\n" + "Options:\n\n" + " -h, --help" + "\t\tShow this help\n" + " -i, --disk" + "\t\tPath of memory unit encrypted with BitLocker\n" + " -o, --outfile" + "\t\tOutput file\n\n", name); + + return EXIT_FAILURE; +} + + +static void fillBuffer(FILE *fp, unsigned char *buffer, int size) +{ + int k; + + for (k = 0; k < size; k++) + buffer[k] = (unsigned char)fgetc(fp); +} + +static void print_hex(unsigned char *str, int len, FILE *out) +{ + int i; + + for (i = 0; i < len; ++i) + fprintf(out, "%02x", str[i]); +} + + +int parse_image(char * encryptedImagePath, char * outputFile) +{ + const char signature[SIGNATURE_LEN] = "-FVE-FS-"; + int version = 0, fileLen = 0, j = 0, i = 0, match = 0; + + unsigned char vmk_entry[4] = { 0x02, 0x00, 0x08, 0x00 }; + unsigned char key_protection_clear[2] = { 0x00, 0x00 }; + unsigned char key_protection_tpm[2] = { 0x00, 0x01 }; + unsigned char key_protection_start_key[2] = { 0x00, 0x02 }; + unsigned char key_protection_recovery[2] = { 0x00, 0x08 }; + unsigned char key_protection_password[2] = { 0x00, 0x20 }; + + unsigned char value_type[2] = { 0x00, 0x05 }; + + unsigned char padding[16] = {0}; + unsigned char tmpLine[128]; + char c,d; + int signature_match=0; + + + printf("Opening file %s\n", encryptedImagePath); + FILE * encryptedImage = fopen(encryptedImagePath, "r"); + if (!encryptedImage) { + fprintf(stderr, "! %s : %s\n", encryptedImagePath, strerror(errno)); + return 1; + } + + fseek(encryptedImage, 0, SEEK_END); + fileLen = ftell(encryptedImage); + fseek(encryptedImage, 0, SEEK_SET); + for (j = 0; j < fileLen; j++) { + c = fgetc(encryptedImage); + while (i < 8 && (unsigned char)c == signature[i]) { + c = fgetc(encryptedImage); + i++; + } + if (i == 8) { + match = 1; + fprintf(stderr, "\nSignature found at 0x%08lx\n", (ftell(encryptedImage) - i - 1)); + fseek(encryptedImage, 1, SEEK_CUR); + version = fgetc(encryptedImage); + fprintf(stderr, "Version: %d ", version); + if (version == 1) + fprintf(stderr, "(Windows Vista)\n"); + else if (version == 2) + fprintf(stderr, "(Windows 7 or later)\n"); + else { + fprintf(stderr, "\nInvalid version, looking for a signature with valid version...\n"); + match = 0; + } + } + if(match == 0) { i=0; continue; } + + i = 0; + while (i < 4 && (unsigned char)c == vmk_entry[i]) { + c = fgetc(encryptedImage); + i++; + } + + if (i == 4) { + fprintf(stderr, "VMK entry found at 0x%08lx\n", (ftell(encryptedImage) - i - 3)); + fseek(encryptedImage, 27, SEEK_CUR); + c = (unsigned char)fgetc(encryptedImage); + d = (unsigned char)fgetc(encryptedImage); + + if ((c == key_protection_clear[0]) && (d == key_protection_clear[1])) + fprintf(stderr, "VMK not encrypted.. stored clear!\n"); + else if ((c == key_protection_tpm[0]) && (d == key_protection_tpm[1])) + fprintf(stderr, "VMK encrypted with TPM...not supported!\n"); + else if ((c == key_protection_start_key[0]) && (d == key_protection_start_key[1])) + fprintf(stderr, "VMK encrypted with Startup Key...not supported!\n"); + else if ((c == key_protection_recovery[0]) && (d == key_protection_recovery[1])) + fprintf(stderr, "VMK encrypted with Recovery key...not supported!\n"); + else if ((c == key_protection_password[0]) && (d == key_protection_password[1])) + { + fprintf(stderr, "VMK encrypted with user password found!\n"); + fseek(encryptedImage, 12, SEEK_CUR); + fillBuffer(encryptedImage, salt, SALT_SIZE); + fseek(encryptedImage, 83, SEEK_CUR); + if (((unsigned char)fgetc(encryptedImage) != value_type[0]) || ((unsigned char)fgetc(encryptedImage) != value_type[1])) { + fprintf(stderr, "Error: VMK not encrypted with AES-CCM\n"); + //ret failure? + } + + fseek(encryptedImage, 3, SEEK_CUR); + fillBuffer(encryptedImage, nonce, NONCE_SIZE); + fillBuffer(encryptedImage, mac, MAC_SIZE); + fillBuffer(encryptedImage, vmk, VMK_SIZE); + break; + } + } + + i = 0; + } + + fclose(encryptedImage); + + if (match == 0) { + fprintf(stderr, "Error while extracting data: No signature found!\n"); + return 1; + } else { + printf("%s result hash:\n$bitlocker$0$%d$", encryptedImagePath, SALT_SIZE); + print_hex(salt, SALT_SIZE, stdout); + printf("$%d$%d$", 0x100000, NONCE_SIZE); // fixed iterations , fixed nonce size + print_hex(nonce, NONCE_SIZE, stdout); + printf("$%d$", VMK_SIZE + 16); + print_hex(padding, 16, stdout); // hack, this should actually be entire AES-CCM encrypted block (which includes vmk) + print_hex(vmk, VMK_SIZE, stdout); + printf("\n"); + + FILE * outFile = fopen(outputFile, "w"); + if (!outFile) { + fprintf(stderr, "! %s : %s\n", outputFile, strerror(errno)); + return 1; + } + + fprintf(outFile, "$bitlocker$0$%d$", SALT_SIZE); + print_hex(salt, SALT_SIZE, outFile); + fprintf(outFile, "$%d$%d$", 0x100000, NONCE_SIZE); // fixed iterations , fixed nonce size + print_hex(nonce, NONCE_SIZE, outFile); + fprintf(outFile, "$%d$", VMK_SIZE + 16); + print_hex(padding, 16, outFile); // hack, this should actually be entire AES-CCM encrypted block (which includes vmk) + print_hex(vmk, VMK_SIZE, outFile); + fprintf(outFile, "\n"); + + fclose(outFile); + } + + return 0; +} + +int main(int argc, char **argv) +{ + errno = 0; + int opt, option_index = 0; + char * imagePath=NULL; + char * outFile=NULL; + + while (1) { + static struct option long_options[] = + { + {"help", no_argument, 0, 'h'}, + {"hashfile", required_argument, 0, 'f'}, + {"dictionary", required_argument, 0, 'd'}, + //{"info", required_argument, 0, 'o'}, + //{"cuda.threads", required_argument, 0, 't'}, + {"passthread", required_argument, 0, 't'}, + {"blocks", required_argument, 0, 'b'}, + {"gpu", required_argument, 0, 'g'}, + {0, 0, 0, 0} + }; + + opt = getopt_long(argc, argv, "hi:o:", long_options, &option_index); + if (opt == -1) + break; + switch (opt) + { + case 'h': + usage(argv[0]); + exit(EXIT_FAILURE); + break; + case 'i': + if(strlen(optarg) >= INPUT_SIZE) + { + fprintf(stderr, "ERROR: Input image path is bigger than %d\n", INPUT_SIZE); + exit(EXIT_FAILURE); + } + imagePath=(char *)Calloc(INPUT_SIZE, sizeof(char)); + strncpy(imagePath, optarg, strlen(optarg)+1); + break; + case 'o': + if(strlen(optarg) >= INPUT_SIZE) + { + fprintf(stderr, "ERROR: Input outfile path is bigger than %d\n", INPUT_SIZE); + exit(EXIT_FAILURE); + } + outFile=(char *)Calloc(INPUT_SIZE, sizeof(char)); + strncpy(outFile,optarg, strlen(optarg)+1); + break; + default: + break; + } + } + + if(!imagePath || !outFile) + { + usage(argv[0]); + exit(EXIT_FAILURE); + } + + printf("\n---------> BitCracker Hash extractor <---------\n"); + parse_image(imagePath, outFile); + + return 0; +} + diff --git a/Hash_Extractor/imgWin10Compatible.txt b/Hash_Extractor/imgWin10Compatible.txt new file mode 100644 index 0000000..2e61aef --- /dev/null +++ b/Hash_Extractor/imgWin10Compatible.txt @@ -0,0 +1 @@ +$bitlocker$0$16$91a4ec232ab95bbb8e8ef964308c6b47$1048576$12$306af9dca50fd30103000000$60$00000000000000000000000000000000509aab04f2161082ed6153d6ea8ad51d45c1ae6ae77cdc470789472640f409a1c2ede715ea5a6bbc320e2312 diff --git a/OpenCL_version/bitcracker.h b/OpenCL_version/bitcracker.h index 445c4e8..bea744a 100755 --- a/OpenCL_version/bitcracker.h +++ b/OpenCL_version/bitcracker.h @@ -1,5 +1,5 @@ /* - * BitCracker: BitLocker password cracking tool, CUDA version. + * BitCracker: BitLocker password cracking tool, OpenCL version. * Copyright (C) 2013-2017 Elena Ago * Massimo Bernaschi * @@ -59,7 +59,7 @@ #define MAC_SIZE 16 #define NONCE_SIZE 12 #define IV_SIZE 16 -#define VMK_SIZE 44 +#define VMK_SIZE 60 #define VMK_DECRYPT_SIZE 16 #define DICT_BUFSIZE (50*1024*1024) #define MAX_PLEN 32 @@ -83,6 +83,10 @@ #define BLOCK_UNIT 32 #define HASH_SIZE_STRING 32 +#define HASH_TAG "$bitlocker$" +#define HASH_TAG_LEN (sizeof(HASH_TAG) - 1) +#define INPUT_HASH_SIZE 210 + #define ATTACK_DEFAULT_THREADS 1024 #define BIT_SUCCESS 0 @@ -183,6 +187,7 @@ extern int platform_id; extern int psw_x_thread; extern int tot_psw; extern size_t size_psw; +extern int strict_check; extern int MAX_PASSWD_SINGLE_KERNEL; extern int DEV_NVIDIA; @@ -201,10 +206,12 @@ extern cl_platform_id cpPlatforms[MAX_NUM_PLATFORMS]; // OpenCL platfo extern cl_uint uiNumDevices; // OpenCL total number of devices extern cl_device_id* cdDevices; // OpenCL device(s) - unsigned int * w_block_precomputed(unsigned char * salt); -int readFilePassword(unsigned char ** buf, int maxNumPsw, FILE *fp); -int readData(char * encryptedImagePath, unsigned char ** salt, unsigned char ** mac, unsigned char ** nonce, unsigned char ** encryptedVMK); +int readFilePassword(char ** buf, int maxNumPsw, FILE *fp); +int parse_data(char *input_hash, + unsigned char ** salt, + unsigned char ** nonce, + unsigned char ** vmk); 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); diff --git a/OpenCL_version/kernel_attack.cl b/OpenCL_version/kernel_attack.cl index 5ecd352..9731911 100755 --- a/OpenCL_version/kernel_attack.cl +++ b/OpenCL_version/kernel_attack.cl @@ -1,7 +1,7 @@ /* * BitCracker: BitLocker password cracking tool, OpenCL version. * Copyright (C) 2013-2017 Elena Ago - * Massimo Bernaschi + * Massimo Bernaschi * * This file is part of BitCracker. * @@ -158,317 +158,6 @@ __constant unsigned int TS3[256] = { 0x4141C382U, 0x9999B029U, 0x2D2D775AU, 0x0F0F111EU, 0xB0B0CB7BU, 0x5454FCA8U, 0xBBBBD66DU, 0x16163A2CU }; -/* -__constant unsigned int TSi0[256] = { - 0x51F4A750U, 0x7E416553U, 0x1A17A4C3U, 0x3A275E96U, 0x3BAB6BCBU, 0x1F9D45F1U, 0xACFA58ABU, 0x4BE30393U, - 0x2030FA55U, 0xAD766DF6U, 0x88CC7691U, 0xF5024C25U, 0x4FE5D7FCU, 0xC52ACBD7U, 0x26354480U, 0xB562A38FU, - 0xDEB15A49U, 0x25BA1B67U, 0x45EA0E98U, 0x5DFEC0E1U, 0xC32F7502U, 0x814CF012U, 0x8D4697A3U, 0x6BD3F9C6U, - 0x038F5FE7U, 0x15929C95U, 0xBF6D7AEBU, 0x955259DAU, 0xD4BE832DU, 0x587421D3U, 0x49E06929U, 0x8EC9C844U, - 0x75C2896AU, 0xF48E7978U, 0x99583E6BU, 0x27B971DDU, 0xBEE14FB6U, 0xF088AD17U, 0xC920AC66U, 0x7DCE3AB4U, - 0x63DF4A18U, 0xE51A3182U, 0x97513360U, 0x62537F45U, 0xB16477E0U, 0xBB6BAE84U, 0xFE81A01CU, 0xF9082B94U, - 0x70486858U, 0x8F45FD19U, 0x94DE6C87U, 0x527BF8B7U, 0xAB73D323U, 0x724B02E2U, 0xE31F8F57U, 0x6655AB2AU, - 0xB2EB2807U, 0x2FB5C203U, 0x86C57B9AU, 0xD33708A5U, 0x302887F2U, 0x23BFA5B2U, 0x02036ABAU, 0xED16825CU, - 0x8ACF1C2BU, 0xA779B492U, 0xF307F2F0U, 0x4E69E2A1U, 0x65DAF4CDU, 0x0605BED5U, 0xD134621FU, 0xC4A6FE8AU, - 0x342E539DU, 0xA2F355A0U, 0x058AE132U, 0xA4F6EB75U, 0x0B83EC39U, 0x4060EFAAU, 0x5E719F06U, 0xBD6E1051U, - 0x3E218AF9U, 0x96DD063DU, 0xDD3E05AEU, 0x4DE6BD46U, 0x91548DB5U, 0x71C45D05U, 0x0406D46FU, 0x605015FFU, - 0x1998FB24U, 0xD6BDE997U, 0x894043CCU, 0x67D99E77U, 0xB0E842BDU, 0x07898B88U, 0xE7195B38U, 0x79C8EEDBU, - 0xA17C0A47U, 0x7C420FE9U, 0xF8841EC9U, 0x00000000U, 0x09808683U, 0x322BED48U, 0x1E1170ACU, 0x6C5A724EU, - 0xFD0EFFFBU, 0x0F853856U, 0x3DAED51EU, 0x362D3927U, 0x0A0FD964U, 0x685CA621U, 0x9B5B54D1U, 0x24362E3AU, - 0x0C0A67B1U, 0x9357E70FU, 0xB4EE96D2U, 0x1B9B919EU, 0x80C0C54FU, 0x61DC20A2U, 0x5A774B69U, 0x1C121A16U, - 0xE293BA0AU, 0xC0A02AE5U, 0x3C22E043U, 0x121B171DU, 0x0E090D0BU, 0xF28BC7ADU, 0x2DB6A8B9U, 0x141EA9C8U, - 0x57F11985U, 0xAF75074CU, 0xEE99DDBBU, 0xA37F60FDU, 0xF701269FU, 0x5C72F5BCU, 0x44663BC5U, 0x5BFB7E34U, - 0x8B432976U, 0xCB23C6DCU, 0xB6EDFC68U, 0xB8E4F163U, 0xD731DCCAU, 0x42638510U, 0x13972240U, 0x84C61120U, - 0x854A247DU, 0xD2BB3DF8U, 0xAEF93211U, 0xC729A16DU, 0x1D9E2F4BU, 0xDCB230F3U, 0x0D8652ECU, 0x77C1E3D0U, - 0x2BB3166CU, 0xA970B999U, 0x119448FAU, 0x47E96422U, 0xA8FC8CC4U, 0xA0F03F1AU, 0x567D2CD8U, 0x223390EFU, - 0x87494EC7U, 0xD938D1C1U, 0x8CCAA2FEU, 0x98D40B36U, 0xA6F581CFU, 0xA57ADE28U, 0xDAB78E26U, 0x3FADBFA4U, - 0x2C3A9DE4U, 0x5078920DU, 0x6A5FCC9BU, 0x547E4662U, 0xF68D13C2U, 0x90D8B8E8U, 0x2E39F75EU, 0x82C3AFF5U, - 0x9F5D80BEU, 0x69D0937CU, 0x6FD52DA9U, 0xCF2512B3U, 0xC8AC993BU, 0x10187DA7U, 0xE89C636EU, 0xDB3BBB7BU, - 0xCD267809U, 0x6E5918F4U, 0xEC9AB701U, 0x834F9AA8U, 0xE6956E65U, 0xAAFFE67EU, 0x21BCCF08U, 0xEF15E8E6U, - 0xBAE79BD9U, 0x4A6F36CEU, 0xEA9F09D4U, 0x29B07CD6U, 0x31A4B2AFU, 0x2A3F2331U, 0xC6A59430U, 0x35A266C0U, - 0x744EBC37U, 0xFC82CAA6U, 0xE090D0B0U, 0x33A7D815U, 0xF104984AU, 0x41ECDAF7U, 0x7FCD500EU, 0x1791F62FU, - 0x764DD68DU, 0x43EFB04DU, 0xCCAA4D54U, 0xE49604DFU, 0x9ED1B5E3U, 0x4C6A881BU, 0xC12C1FB8U, 0x4665517FU, - 0x9D5EEA04U, 0x018C355DU, 0xFA877473U, 0xFB0B412EU, 0xB3671D5AU, 0x92DBD252U, 0xE9105633U, 0x6DD64713U, - 0x9AD7618CU, 0x37A10C7AU, 0x59F8148EU, 0xEB133C89U, 0xCEA927EEU, 0xB761C935U, 0xE11CE5EDU, 0x7A47B13CU, - 0x9CD2DF59U, 0x55F2733FU, 0x1814CE79U, 0x73C737BFU, 0x53F7CDEAU, 0x5FFDAA5BU, 0xDF3D6F14U, 0x7844DB86U, - 0xCAAFF381U, 0xB968C43EU, 0x3824342CU, 0xC2A3405FU, 0x161DC372U, 0xBCE2250CU, 0x283C498BU, 0xFF0D9541U, - 0x39A80171U, 0x080CB3DEU, 0xD8B4E49CU, 0x6456C190U, 0x7BCB8461U, 0xD532B670U, 0x486C5C74U, 0xD0B85742U -}; - -__constant unsigned int TSi1[256] = { - 0x5051F4A7U, 0x537E4165U, 0xC31A17A4U, 0x963A275EU, 0xCB3BAB6BU, 0xF11F9D45U, 0xABACFA58U, 0x934BE303U, - 0x552030FAU, 0xF6AD766DU, 0x9188CC76U, 0x25F5024CU, 0xFC4FE5D7U, 0xD7C52ACBU, 0x80263544U, 0x8FB562A3U, - 0x49DEB15AU, 0x6725BA1BU, 0x9845EA0EU, 0xE15DFEC0U, 0x02C32F75U, 0x12814CF0U, 0xA38D4697U, 0xC66BD3F9U, - 0xE7038F5FU, 0x9515929CU, 0xEBBF6D7AU, 0xDA955259U, 0x2DD4BE83U, 0xD3587421U, 0x2949E069U, 0x448EC9C8U, - 0x6A75C289U, 0x78F48E79U, 0x6B99583EU, 0xDD27B971U, 0xB6BEE14FU, 0x17F088ADU, 0x66C920ACU, 0xB47DCE3AU, - 0x1863DF4AU, 0x82E51A31U, 0x60975133U, 0x4562537FU, 0xE0B16477U, 0x84BB6BAEU, 0x1CFE81A0U, 0x94F9082BU, - 0x58704868U, 0x198F45FDU, 0x8794DE6CU, 0xB7527BF8U, 0x23AB73D3U, 0xE2724B02U, 0x57E31F8FU, 0x2A6655ABU, - 0x07B2EB28U, 0x032FB5C2U, 0x9A86C57BU, 0xA5D33708U, 0xF2302887U, 0xB223BFA5U, 0xBA02036AU, 0x5CED1682U, - 0x2B8ACF1CU, 0x92A779B4U, 0xF0F307F2U, 0xA14E69E2U, 0xCD65DAF4U, 0xD50605BEU, 0x1FD13462U, 0x8AC4A6FEU, - 0x9D342E53U, 0xA0A2F355U, 0x32058AE1U, 0x75A4F6EBU, 0x390B83ECU, 0xAA4060EFU, 0x065E719FU, 0x51BD6E10U, - 0xF93E218AU, 0x3D96DD06U, 0xAEDD3E05U, 0x464DE6BDU, 0xB591548DU, 0x0571C45DU, 0x6F0406D4U, 0xFF605015U, - 0x241998FBU, 0x97D6BDE9U, 0xCC894043U, 0x7767D99EU, 0xBDB0E842U, 0x8807898BU, 0x38E7195BU, 0xDB79C8EEU, - 0x47A17C0AU, 0xE97C420FU, 0xC9F8841EU, 0x00000000U, 0x83098086U, 0x48322BEDU, 0xAC1E1170U, 0x4E6C5A72U, - 0xFBFD0EFFU, 0x560F8538U, 0x1E3DAED5U, 0x27362D39U, 0x640A0FD9U, 0x21685CA6U, 0xD19B5B54U, 0x3A24362EU, - 0xB10C0A67U, 0x0F9357E7U, 0xD2B4EE96U, 0x9E1B9B91U, 0x4F80C0C5U, 0xA261DC20U, 0x695A774BU, 0x161C121AU, - 0x0AE293BAU, 0xE5C0A02AU, 0x433C22E0U, 0x1D121B17U, 0x0B0E090DU, 0xADF28BC7U, 0xB92DB6A8U, 0xC8141EA9U, - 0x8557F119U, 0x4CAF7507U, 0xBBEE99DDU, 0xFDA37F60U, 0x9FF70126U, 0xBC5C72F5U, 0xC544663BU, 0x345BFB7EU, - 0x768B4329U, 0xDCCB23C6U, 0x68B6EDFCU, 0x63B8E4F1U, 0xCAD731DCU, 0x10426385U, 0x40139722U, 0x2084C611U, - 0x7D854A24U, 0xF8D2BB3DU, 0x11AEF932U, 0x6DC729A1U, 0x4B1D9E2FU, 0xF3DCB230U, 0xEC0D8652U, 0xD077C1E3U, - 0x6C2BB316U, 0x99A970B9U, 0xFA119448U, 0x2247E964U, 0xC4A8FC8CU, 0x1AA0F03FU, 0xD8567D2CU, 0xEF223390U, - 0xC787494EU, 0xC1D938D1U, 0xFE8CCAA2U, 0x3698D40BU, 0xCFA6F581U, 0x28A57ADEU, 0x26DAB78EU, 0xA43FADBFU, - 0xE42C3A9DU, 0x0D507892U, 0x9B6A5FCCU, 0x62547E46U, 0xC2F68D13U, 0xE890D8B8U, 0x5E2E39F7U, 0xF582C3AFU, - 0xBE9F5D80U, 0x7C69D093U, 0xA96FD52DU, 0xB3CF2512U, 0x3BC8AC99U, 0xA710187DU, 0x6EE89C63U, 0x7BDB3BBBU, - 0x09CD2678U, 0xF46E5918U, 0x01EC9AB7U, 0xA8834F9AU, 0x65E6956EU, 0x7EAAFFE6U, 0x0821BCCFU, 0xE6EF15E8U, - 0xD9BAE79BU, 0xCE4A6F36U, 0xD4EA9F09U, 0xD629B07CU, 0xAF31A4B2U, 0x312A3F23U, 0x30C6A594U, 0xC035A266U, - 0x37744EBCU, 0xA6FC82CAU, 0xB0E090D0U, 0x1533A7D8U, 0x4AF10498U, 0xF741ECDAU, 0x0E7FCD50U, 0x2F1791F6U, - 0x8D764DD6U, 0x4D43EFB0U, 0x54CCAA4DU, 0xDFE49604U, 0xE39ED1B5U, 0x1B4C6A88U, 0xB8C12C1FU, 0x7F466551U, - 0x049D5EEAU, 0x5D018C35U, 0x73FA8774U, 0x2EFB0B41U, 0x5AB3671DU, 0x5292DBD2U, 0x33E91056U, 0x136DD647U, - 0x8C9AD761U, 0x7A37A10CU, 0x8E59F814U, 0x89EB133CU, 0xEECEA927U, 0x35B761C9U, 0xEDE11CE5U, 0x3C7A47B1U, - 0x599CD2DFU, 0x3F55F273U, 0x791814CEU, 0xBF73C737U, 0xEA53F7CDU, 0x5B5FFDAAU, 0x14DF3D6FU, 0x867844DBU, - 0x81CAAFF3U, 0x3EB968C4U, 0x2C382434U, 0x5FC2A340U, 0x72161DC3U, 0x0CBCE225U, 0x8B283C49U, 0x41FF0D95U, - 0x7139A801U, 0xDE080CB3U, 0x9CD8B4E4U, 0x906456C1U, 0x617BCB84U, 0x70D532B6U, 0x74486C5CU, 0x42D0B857U -}; - -__constant unsigned int TSi2[256] = { - 0xA75051F4U, 0x65537E41U, 0xA4C31A17U, 0x5E963A27U, 0x6BCB3BABU, 0x45F11F9DU, 0x58ABACFAU, 0x03934BE3U, - 0xFA552030U, 0x6DF6AD76U, 0x769188CCU, 0x4C25F502U, 0xD7FC4FE5U, 0xCBD7C52AU, 0x44802635U, 0xA38FB562U, - 0x5A49DEB1U, 0x1B6725BAU, 0x0E9845EAU, 0xC0E15DFEU, 0x7502C32FU, 0xF012814CU, 0x97A38D46U, 0xF9C66BD3U, - 0x5FE7038FU, 0x9C951592U, 0x7AEBBF6DU, 0x59DA9552U, 0x832DD4BEU, 0x21D35874U, 0x692949E0U, 0xC8448EC9U, - 0x896A75C2U, 0x7978F48EU, 0x3E6B9958U, 0x71DD27B9U, 0x4FB6BEE1U, 0xAD17F088U, 0xAC66C920U, 0x3AB47DCEU, - 0x4A1863DFU, 0x3182E51AU, 0x33609751U, 0x7F456253U, 0x77E0B164U, 0xAE84BB6BU, 0xA01CFE81U, 0x2B94F908U, - 0x68587048U, 0xFD198F45U, 0x6C8794DEU, 0xF8B7527BU, 0xD323AB73U, 0x02E2724BU, 0x8F57E31FU, 0xAB2A6655U, - 0x2807B2EBU, 0xC2032FB5U, 0x7B9A86C5U, 0x08A5D337U, 0x87F23028U, 0xA5B223BFU, 0x6ABA0203U, 0x825CED16U, - 0x1C2B8ACFU, 0xB492A779U, 0xF2F0F307U, 0xE2A14E69U, 0xF4CD65DAU, 0xBED50605U, 0x621FD134U, 0xFE8AC4A6U, - 0x539D342EU, 0x55A0A2F3U, 0xE132058AU, 0xEB75A4F6U, 0xEC390B83U, 0xEFAA4060U, 0x9F065E71U, 0x1051BD6EU, - 0x8AF93E21U, 0x063D96DDU, 0x05AEDD3EU, 0xBD464DE6U, 0x8DB59154U, 0x5D0571C4U, 0xD46F0406U, 0x15FF6050U, - 0xFB241998U, 0xE997D6BDU, 0x43CC8940U, 0x9E7767D9U, 0x42BDB0E8U, 0x8B880789U, 0x5B38E719U, 0xEEDB79C8U, - 0x0A47A17CU, 0x0FE97C42U, 0x1EC9F884U, 0x00000000U, 0x86830980U, 0xED48322BU, 0x70AC1E11U, 0x724E6C5AU, - 0xFFFBFD0EU, 0x38560F85U, 0xD51E3DAEU, 0x3927362DU, 0xD9640A0FU, 0xA621685CU, 0x54D19B5BU, 0x2E3A2436U, - 0x67B10C0AU, 0xE70F9357U, 0x96D2B4EEU, 0x919E1B9BU, 0xC54F80C0U, 0x20A261DCU, 0x4B695A77U, 0x1A161C12U, - 0xBA0AE293U, 0x2AE5C0A0U, 0xE0433C22U, 0x171D121BU, 0x0D0B0E09U, 0xC7ADF28BU, 0xA8B92DB6U, 0xA9C8141EU, - 0x198557F1U, 0x074CAF75U, 0xDDBBEE99U, 0x60FDA37FU, 0x269FF701U, 0xF5BC5C72U, 0x3BC54466U, 0x7E345BFBU, - 0x29768B43U, 0xC6DCCB23U, 0xFC68B6EDU, 0xF163B8E4U, 0xDCCAD731U, 0x85104263U, 0x22401397U, 0x112084C6U, - 0x247D854AU, 0x3DF8D2BBU, 0x3211AEF9U, 0xA16DC729U, 0x2F4B1D9EU, 0x30F3DCB2U, 0x52EC0D86U, 0xE3D077C1U, - 0x166C2BB3U, 0xB999A970U, 0x48FA1194U, 0x642247E9U, 0x8CC4A8FCU, 0x3F1AA0F0U, 0x2CD8567DU, 0x90EF2233U, - 0x4EC78749U, 0xD1C1D938U, 0xA2FE8CCAU, 0x0B3698D4U, 0x81CFA6F5U, 0xDE28A57AU, 0x8E26DAB7U, 0xBFA43FADU, - 0x9DE42C3AU, 0x920D5078U, 0xCC9B6A5FU, 0x4662547EU, 0x13C2F68DU, 0xB8E890D8U, 0xF75E2E39U, 0xAFF582C3U, - 0x80BE9F5DU, 0x937C69D0U, 0x2DA96FD5U, 0x12B3CF25U, 0x993BC8ACU, 0x7DA71018U, 0x636EE89CU, 0xBB7BDB3BU, - 0x7809CD26U, 0x18F46E59U, 0xB701EC9AU, 0x9AA8834FU, 0x6E65E695U, 0xE67EAAFFU, 0xCF0821BCU, 0xE8E6EF15U, - 0x9BD9BAE7U, 0x36CE4A6FU, 0x09D4EA9FU, 0x7CD629B0U, 0xB2AF31A4U, 0x23312A3FU, 0x9430C6A5U, 0x66C035A2U, - 0xBC37744EU, 0xCAA6FC82U, 0xD0B0E090U, 0xD81533A7U, 0x984AF104U, 0xDAF741ECU, 0x500E7FCDU, 0xF62F1791U, - 0xD68D764DU, 0xB04D43EFU, 0x4D54CCAAU, 0x04DFE496U, 0xB5E39ED1U, 0x881B4C6AU, 0x1FB8C12CU, 0x517F4665U, - 0xEA049D5EU, 0x355D018CU, 0x7473FA87U, 0x412EFB0BU, 0x1D5AB367U, 0xD25292DBU, 0x5633E910U, 0x47136DD6U, - 0x618C9AD7U, 0x0C7A37A1U, 0x148E59F8U, 0x3C89EB13U, 0x27EECEA9U, 0xC935B761U, 0xE5EDE11CU, 0xB13C7A47U, - 0xDF599CD2U, 0x733F55F2U, 0xCE791814U, 0x37BF73C7U, 0xCDEA53F7U, 0xAA5B5FFDU, 0x6F14DF3DU, 0xDB867844U, - 0xF381CAAFU, 0xC43EB968U, 0x342C3824U, 0x405FC2A3U, 0xC372161DU, 0x250CBCE2U, 0x498B283CU, 0x9541FF0DU, - 0x017139A8U, 0xB3DE080CU, 0xE49CD8B4U, 0xC1906456U, 0x84617BCBU, 0xB670D532U, 0x5C74486CU, 0x5742D0B8U -}; - -__constant unsigned int TSi3[256] = { - 0xF4A75051U, 0x4165537EU, 0x17A4C31AU, 0x275E963AU, 0xAB6BCB3BU, 0x9D45F11FU, 0xFA58ABACU, 0xE303934BU, - 0x30FA5520U, 0x766DF6ADU, 0xCC769188U, 0x024C25F5U, 0xE5D7FC4FU, 0x2ACBD7C5U, 0x35448026U, 0x62A38FB5U, - 0xB15A49DEU, 0xBA1B6725U, 0xEA0E9845U, 0xFEC0E15DU, 0x2F7502C3U, 0x4CF01281U, 0x4697A38DU, 0xD3F9C66BU, - 0x8F5FE703U, 0x929C9515U, 0x6D7AEBBFU, 0x5259DA95U, 0xBE832DD4U, 0x7421D358U, 0xE0692949U, 0xC9C8448EU, - 0xC2896A75U, 0x8E7978F4U, 0x583E6B99U, 0xB971DD27U, 0xE14FB6BEU, 0x88AD17F0U, 0x20AC66C9U, 0xCE3AB47DU, - 0xDF4A1863U, 0x1A3182E5U, 0x51336097U, 0x537F4562U, 0x6477E0B1U, 0x6BAE84BBU, 0x81A01CFEU, 0x082B94F9U, - 0x48685870U, 0x45FD198FU, 0xDE6C8794U, 0x7BF8B752U, 0x73D323ABU, 0x4B02E272U, 0x1F8F57E3U, 0x55AB2A66U, - 0xEB2807B2U, 0xB5C2032FU, 0xC57B9A86U, 0x3708A5D3U, 0x2887F230U, 0xBFA5B223U, 0x036ABA02U, 0x16825CEDU, - 0xCF1C2B8AU, 0x79B492A7U, 0x07F2F0F3U, 0x69E2A14EU, 0xDAF4CD65U, 0x05BED506U, 0x34621FD1U, 0xA6FE8AC4U, - 0x2E539D34U, 0xF355A0A2U, 0x8AE13205U, 0xF6EB75A4U, 0x83EC390BU, 0x60EFAA40U, 0x719F065EU, 0x6E1051BDU, - 0x218AF93EU, 0xDD063D96U, 0x3E05AEDDU, 0xE6BD464DU, 0x548DB591U, 0xC45D0571U, 0x06D46F04U, 0x5015FF60U, - 0x98FB2419U, 0xBDE997D6U, 0x4043CC89U, 0xD99E7767U, 0xE842BDB0U, 0x898B8807U, 0x195B38E7U, 0xC8EEDB79U, - 0x7C0A47A1U, 0x420FE97CU, 0x841EC9F8U, 0x00000000U, 0x80868309U, 0x2BED4832U, 0x1170AC1EU, 0x5A724E6CU, - 0x0EFFFBFDU, 0x8538560FU, 0xAED51E3DU, 0x2D392736U, 0x0FD9640AU, 0x5CA62168U, 0x5B54D19BU, 0x362E3A24U, - 0x0A67B10CU, 0x57E70F93U, 0xEE96D2B4U, 0x9B919E1BU, 0xC0C54F80U, 0xDC20A261U, 0x774B695AU, 0x121A161CU, - 0x93BA0AE2U, 0xA02AE5C0U, 0x22E0433CU, 0x1B171D12U, 0x090D0B0EU, 0x8BC7ADF2U, 0xB6A8B92DU, 0x1EA9C814U, - 0xF1198557U, 0x75074CAFU, 0x99DDBBEEU, 0x7F60FDA3U, 0x01269FF7U, 0x72F5BC5CU, 0x663BC544U, 0xFB7E345BU, - 0x4329768BU, 0x23C6DCCBU, 0xEDFC68B6U, 0xE4F163B8U, 0x31DCCAD7U, 0x63851042U, 0x97224013U, 0xC6112084U, - 0x4A247D85U, 0xBB3DF8D2U, 0xF93211AEU, 0x29A16DC7U, 0x9E2F4B1DU, 0xB230F3DCU, 0x8652EC0DU, 0xC1E3D077U, - 0xB3166C2BU, 0x70B999A9U, 0x9448FA11U, 0xE9642247U, 0xFC8CC4A8U, 0xF03F1AA0U, 0x7D2CD856U, 0x3390EF22U, - 0x494EC787U, 0x38D1C1D9U, 0xCAA2FE8CU, 0xD40B3698U, 0xF581CFA6U, 0x7ADE28A5U, 0xB78E26DAU, 0xADBFA43FU, - 0x3A9DE42CU, 0x78920D50U, 0x5FCC9B6AU, 0x7E466254U, 0x8D13C2F6U, 0xD8B8E890U, 0x39F75E2EU, 0xC3AFF582U, - 0x5D80BE9FU, 0xD0937C69U, 0xD52DA96FU, 0x2512B3CFU, 0xAC993BC8U, 0x187DA710U, 0x9C636EE8U, 0x3BBB7BDBU, - 0x267809CDU, 0x5918F46EU, 0x9AB701ECU, 0x4F9AA883U, 0x956E65E6U, 0xFFE67EAAU, 0xBCCF0821U, 0x15E8E6EFU, - 0xE79BD9BAU, 0x6F36CE4AU, 0x9F09D4EAU, 0xB07CD629U, 0xA4B2AF31U, 0x3F23312AU, 0xA59430C6U, 0xA266C035U, - 0x4EBC3774U, 0x82CAA6FCU, 0x90D0B0E0U, 0xA7D81533U, 0x04984AF1U, 0xECDAF741U, 0xCD500E7FU, 0x91F62F17U, - 0x4DD68D76U, 0xEFB04D43U, 0xAA4D54CCU, 0x9604DFE4U, 0xD1B5E39EU, 0x6A881B4CU, 0x2C1FB8C1U, 0x65517F46U, - 0x5EEA049DU, 0x8C355D01U, 0x877473FAU, 0x0B412EFBU, 0x671D5AB3U, 0xDBD25292U, 0x105633E9U, 0xD647136DU, - 0xD7618C9AU, 0xA10C7A37U, 0xF8148E59U, 0x133C89EBU, 0xA927EECEU, 0x61C935B7U, 0x1CE5EDE1U, 0x47B13C7AU, - 0xD2DF599CU, 0xF2733F55U, 0x14CE7918U, 0xC737BF73U, 0xF7CDEA53U, 0xFDAA5B5FU, 0x3D6F14DFU, 0x44DB8678U, - 0xAFF381CAU, 0x68C43EB9U, 0x24342C38U, 0xA3405FC2U, 0x1DC37216U, 0xE2250CBCU, 0x3C498B28U, 0x0D9541FFU, - 0xA8017139U, 0x0CB3DE08U, 0xB4E49CD8U, 0x56C19064U, 0xCB84617BU, 0x32B670D5U, 0x6C5C7448U, 0xB85742D0U -}; -*/ -/* -// these are the InvMixColumn wthiout the SBOXI[] -__constant unsigned int USi0[256] = { - 0x00000000U, 0x0E090D0BU, 0x1C121A16U, 0x121B171DU, 0x3824342CU, 0x362D3927U, 0x24362E3AU, 0x2A3F2331U, - 0x70486858U, 0x7E416553U, 0x6C5A724EU, 0x62537F45U, 0x486C5C74U, 0x4665517FU, 0x547E4662U, 0x5A774B69U, - 0xE090D0B0U, 0xEE99DDBBU, 0xFC82CAA6U, 0xF28BC7ADU, 0xD8B4E49CU, 0xD6BDE997U, 0xC4A6FE8AU, 0xCAAFF381U, - 0x90D8B8E8U, 0x9ED1B5E3U, 0x8CCAA2FEU, 0x82C3AFF5U, 0xA8FC8CC4U, 0xA6F581CFU, 0xB4EE96D2U, 0xBAE79BD9U, - 0xDB3BBB7BU, 0xD532B670U, 0xC729A16DU, 0xC920AC66U, 0xE31F8F57U, 0xED16825CU, 0xFF0D9541U, 0xF104984AU, - 0xAB73D323U, 0xA57ADE28U, 0xB761C935U, 0xB968C43EU, 0x9357E70FU, 0x9D5EEA04U, 0x8F45FD19U, 0x814CF012U, - 0x3BAB6BCBU, 0x35A266C0U, 0x27B971DDU, 0x29B07CD6U, 0x038F5FE7U, 0x0D8652ECU, 0x1F9D45F1U, 0x119448FAU, - 0x4BE30393U, 0x45EA0E98U, 0x57F11985U, 0x59F8148EU, 0x73C737BFU, 0x7DCE3AB4U, 0x6FD52DA9U, 0x61DC20A2U, - 0xAD766DF6U, 0xA37F60FDU, 0xB16477E0U, 0xBF6D7AEBU, 0x955259DAU, 0x9B5B54D1U, 0x894043CCU, 0x87494EC7U, - 0xDD3E05AEU, 0xD33708A5U, 0xC12C1FB8U, 0xCF2512B3U, 0xE51A3182U, 0xEB133C89U, 0xF9082B94U, 0xF701269FU, - 0x4DE6BD46U, 0x43EFB04DU, 0x51F4A750U, 0x5FFDAA5BU, 0x75C2896AU, 0x7BCB8461U, 0x69D0937CU, 0x67D99E77U, - 0x3DAED51EU, 0x33A7D815U, 0x21BCCF08U, 0x2FB5C203U, 0x058AE132U, 0x0B83EC39U, 0x1998FB24U, 0x1791F62FU, - 0x764DD68DU, 0x7844DB86U, 0x6A5FCC9BU, 0x6456C190U, 0x4E69E2A1U, 0x4060EFAAU, 0x527BF8B7U, 0x5C72F5BCU, - 0x0605BED5U, 0x080CB3DEU, 0x1A17A4C3U, 0x141EA9C8U, 0x3E218AF9U, 0x302887F2U, 0x223390EFU, 0x2C3A9DE4U, - 0x96DD063DU, 0x98D40B36U, 0x8ACF1C2BU, 0x84C61120U, 0xAEF93211U, 0xA0F03F1AU, 0xB2EB2807U, 0xBCE2250CU, - 0xE6956E65U, 0xE89C636EU, 0xFA877473U, 0xF48E7978U, 0xDEB15A49U, 0xD0B85742U, 0xC2A3405FU, 0xCCAA4D54U, - 0x41ECDAF7U, 0x4FE5D7FCU, 0x5DFEC0E1U, 0x53F7CDEAU, 0x79C8EEDBU, 0x77C1E3D0U, 0x65DAF4CDU, 0x6BD3F9C6U, - 0x31A4B2AFU, 0x3FADBFA4U, 0x2DB6A8B9U, 0x23BFA5B2U, 0x09808683U, 0x07898B88U, 0x15929C95U, 0x1B9B919EU, - 0xA17C0A47U, 0xAF75074CU, 0xBD6E1051U, 0xB3671D5AU, 0x99583E6BU, 0x97513360U, 0x854A247DU, 0x8B432976U, - 0xD134621FU, 0xDF3D6F14U, 0xCD267809U, 0xC32F7502U, 0xE9105633U, 0xE7195B38U, 0xF5024C25U, 0xFB0B412EU, - 0x9AD7618CU, 0x94DE6C87U, 0x86C57B9AU, 0x88CC7691U, 0xA2F355A0U, 0xACFA58ABU, 0xBEE14FB6U, 0xB0E842BDU, - 0xEA9F09D4U, 0xE49604DFU, 0xF68D13C2U, 0xF8841EC9U, 0xD2BB3DF8U, 0xDCB230F3U, 0xCEA927EEU, 0xC0A02AE5U, - 0x7A47B13CU, 0x744EBC37U, 0x6655AB2AU, 0x685CA621U, 0x42638510U, 0x4C6A881BU, 0x5E719F06U, 0x5078920DU, - 0x0A0FD964U, 0x0406D46FU, 0x161DC372U, 0x1814CE79U, 0x322BED48U, 0x3C22E043U, 0x2E39F75EU, 0x2030FA55U, - 0xEC9AB701U, 0xE293BA0AU, 0xF088AD17U, 0xFE81A01CU, 0xD4BE832DU, 0xDAB78E26U, 0xC8AC993BU, 0xC6A59430U, - 0x9CD2DF59U, 0x92DBD252U, 0x80C0C54FU, 0x8EC9C844U, 0xA4F6EB75U, 0xAAFFE67EU, 0xB8E4F163U, 0xB6EDFC68U, - 0x0C0A67B1U, 0x02036ABAU, 0x10187DA7U, 0x1E1170ACU, 0x342E539DU, 0x3A275E96U, 0x283C498BU, 0x26354480U, - 0x7C420FE9U, 0x724B02E2U, 0x605015FFU, 0x6E5918F4U, 0x44663BC5U, 0x4A6F36CEU, 0x587421D3U, 0x567D2CD8U, - 0x37A10C7AU, 0x39A80171U, 0x2BB3166CU, 0x25BA1B67U, 0x0F853856U, 0x018C355DU, 0x13972240U, 0x1D9E2F4BU, - 0x47E96422U, 0x49E06929U, 0x5BFB7E34U, 0x55F2733FU, 0x7FCD500EU, 0x71C45D05U, 0x63DF4A18U, 0x6DD64713U, - 0xD731DCCAU, 0xD938D1C1U, 0xCB23C6DCU, 0xC52ACBD7U, 0xEF15E8E6U, 0xE11CE5EDU, 0xF307F2F0U, 0xFD0EFFFBU, - 0xA779B492U, 0xA970B999U, 0xBB6BAE84U, 0xB562A38FU, 0x9F5D80BEU, 0x91548DB5U, 0x834F9AA8U, 0x8D4697A3U -}; - -__constant unsigned int USi1[256] = { - 0x00000000U, 0x0B0E090DU, 0x161C121AU, 0x1D121B17U, 0x2C382434U, 0x27362D39U, 0x3A24362EU, 0x312A3F23U, - 0x58704868U, 0x537E4165U, 0x4E6C5A72U, 0x4562537FU, 0x74486C5CU, 0x7F466551U, 0x62547E46U, 0x695A774BU, - 0xB0E090D0U, 0xBBEE99DDU, 0xA6FC82CAU, 0xADF28BC7U, 0x9CD8B4E4U, 0x97D6BDE9U, 0x8AC4A6FEU, 0x81CAAFF3U, - 0xE890D8B8U, 0xE39ED1B5U, 0xFE8CCAA2U, 0xF582C3AFU, 0xC4A8FC8CU, 0xCFA6F581U, 0xD2B4EE96U, 0xD9BAE79BU, - 0x7BDB3BBBU, 0x70D532B6U, 0x6DC729A1U, 0x66C920ACU, 0x57E31F8FU, 0x5CED1682U, 0x41FF0D95U, 0x4AF10498U, - 0x23AB73D3U, 0x28A57ADEU, 0x35B761C9U, 0x3EB968C4U, 0x0F9357E7U, 0x049D5EEAU, 0x198F45FDU, 0x12814CF0U, - 0xCB3BAB6BU, 0xC035A266U, 0xDD27B971U, 0xD629B07CU, 0xE7038F5FU, 0xEC0D8652U, 0xF11F9D45U, 0xFA119448U, - 0x934BE303U, 0x9845EA0EU, 0x8557F119U, 0x8E59F814U, 0xBF73C737U, 0xB47DCE3AU, 0xA96FD52DU, 0xA261DC20U, - 0xF6AD766DU, 0xFDA37F60U, 0xE0B16477U, 0xEBBF6D7AU, 0xDA955259U, 0xD19B5B54U, 0xCC894043U, 0xC787494EU, - 0xAEDD3E05U, 0xA5D33708U, 0xB8C12C1FU, 0xB3CF2512U, 0x82E51A31U, 0x89EB133CU, 0x94F9082BU, 0x9FF70126U, - 0x464DE6BDU, 0x4D43EFB0U, 0x5051F4A7U, 0x5B5FFDAAU, 0x6A75C289U, 0x617BCB84U, 0x7C69D093U, 0x7767D99EU, - 0x1E3DAED5U, 0x1533A7D8U, 0x0821BCCFU, 0x032FB5C2U, 0x32058AE1U, 0x390B83ECU, 0x241998FBU, 0x2F1791F6U, - 0x8D764DD6U, 0x867844DBU, 0x9B6A5FCCU, 0x906456C1U, 0xA14E69E2U, 0xAA4060EFU, 0xB7527BF8U, 0xBC5C72F5U, - 0xD50605BEU, 0xDE080CB3U, 0xC31A17A4U, 0xC8141EA9U, 0xF93E218AU, 0xF2302887U, 0xEF223390U, 0xE42C3A9DU, - 0x3D96DD06U, 0x3698D40BU, 0x2B8ACF1CU, 0x2084C611U, 0x11AEF932U, 0x1AA0F03FU, 0x07B2EB28U, 0x0CBCE225U, - 0x65E6956EU, 0x6EE89C63U, 0x73FA8774U, 0x78F48E79U, 0x49DEB15AU, 0x42D0B857U, 0x5FC2A340U, 0x54CCAA4DU, - 0xF741ECDAU, 0xFC4FE5D7U, 0xE15DFEC0U, 0xEA53F7CDU, 0xDB79C8EEU, 0xD077C1E3U, 0xCD65DAF4U, 0xC66BD3F9U, - 0xAF31A4B2U, 0xA43FADBFU, 0xB92DB6A8U, 0xB223BFA5U, 0x83098086U, 0x8807898BU, 0x9515929CU, 0x9E1B9B91U, - 0x47A17C0AU, 0x4CAF7507U, 0x51BD6E10U, 0x5AB3671DU, 0x6B99583EU, 0x60975133U, 0x7D854A24U, 0x768B4329U, - 0x1FD13462U, 0x14DF3D6FU, 0x09CD2678U, 0x02C32F75U, 0x33E91056U, 0x38E7195BU, 0x25F5024CU, 0x2EFB0B41U, - 0x8C9AD761U, 0x8794DE6CU, 0x9A86C57BU, 0x9188CC76U, 0xA0A2F355U, 0xABACFA58U, 0xB6BEE14FU, 0xBDB0E842U, - 0xD4EA9F09U, 0xDFE49604U, 0xC2F68D13U, 0xC9F8841EU, 0xF8D2BB3DU, 0xF3DCB230U, 0xEECEA927U, 0xE5C0A02AU, - 0x3C7A47B1U, 0x37744EBCU, 0x2A6655ABU, 0x21685CA6U, 0x10426385U, 0x1B4C6A88U, 0x065E719FU, 0x0D507892U, - 0x640A0FD9U, 0x6F0406D4U, 0x72161DC3U, 0x791814CEU, 0x48322BEDU, 0x433C22E0U, 0x5E2E39F7U, 0x552030FAU, - 0x01EC9AB7U, 0x0AE293BAU, 0x17F088ADU, 0x1CFE81A0U, 0x2DD4BE83U, 0x26DAB78EU, 0x3BC8AC99U, 0x30C6A594U, - 0x599CD2DFU, 0x5292DBD2U, 0x4F80C0C5U, 0x448EC9C8U, 0x75A4F6EBU, 0x7EAAFFE6U, 0x63B8E4F1U, 0x68B6EDFCU, - 0xB10C0A67U, 0xBA02036AU, 0xA710187DU, 0xAC1E1170U, 0x9D342E53U, 0x963A275EU, 0x8B283C49U, 0x80263544U, - 0xE97C420FU, 0xE2724B02U, 0xFF605015U, 0xF46E5918U, 0xC544663BU, 0xCE4A6F36U, 0xD3587421U, 0xD8567D2CU, - 0x7A37A10CU, 0x7139A801U, 0x6C2BB316U, 0x6725BA1BU, 0x560F8538U, 0x5D018C35U, 0x40139722U, 0x4B1D9E2FU, - 0x2247E964U, 0x2949E069U, 0x345BFB7EU, 0x3F55F273U, 0x0E7FCD50U, 0x0571C45DU, 0x1863DF4AU, 0x136DD647U, - 0xCAD731DCU, 0xC1D938D1U, 0xDCCB23C6U, 0xD7C52ACBU, 0xE6EF15E8U, 0xEDE11CE5U, 0xF0F307F2U, 0xFBFD0EFFU, - 0x92A779B4U, 0x99A970B9U, 0x84BB6BAEU, 0x8FB562A3U, 0xBE9F5D80U, 0xB591548DU, 0xA8834F9AU, 0xA38D4697U -}; - -__constant unsigned int USi2[256] = { - 0x00000000U, 0x0D0B0E09U, 0x1A161C12U, 0x171D121BU, 0x342C3824U, 0x3927362DU, 0x2E3A2436U, 0x23312A3FU, - 0x68587048U, 0x65537E41U, 0x724E6C5AU, 0x7F456253U, 0x5C74486CU, 0x517F4665U, 0x4662547EU, 0x4B695A77U, - 0xD0B0E090U, 0xDDBBEE99U, 0xCAA6FC82U, 0xC7ADF28BU, 0xE49CD8B4U, 0xE997D6BDU, 0xFE8AC4A6U, 0xF381CAAFU, - 0xB8E890D8U, 0xB5E39ED1U, 0xA2FE8CCAU, 0xAFF582C3U, 0x8CC4A8FCU, 0x81CFA6F5U, 0x96D2B4EEU, 0x9BD9BAE7U, - 0xBB7BDB3BU, 0xB670D532U, 0xA16DC729U, 0xAC66C920U, 0x8F57E31FU, 0x825CED16U, 0x9541FF0DU, 0x984AF104U, - 0xD323AB73U, 0xDE28A57AU, 0xC935B761U, 0xC43EB968U, 0xE70F9357U, 0xEA049D5EU, 0xFD198F45U, 0xF012814CU, - 0x6BCB3BABU, 0x66C035A2U, 0x71DD27B9U, 0x7CD629B0U, 0x5FE7038FU, 0x52EC0D86U, 0x45F11F9DU, 0x48FA1194U, - 0x03934BE3U, 0x0E9845EAU, 0x198557F1U, 0x148E59F8U, 0x37BF73C7U, 0x3AB47DCEU, 0x2DA96FD5U, 0x20A261DCU, - 0x6DF6AD76U, 0x60FDA37FU, 0x77E0B164U, 0x7AEBBF6DU, 0x59DA9552U, 0x54D19B5BU, 0x43CC8940U, 0x4EC78749U, - 0x05AEDD3EU, 0x08A5D337U, 0x1FB8C12CU, 0x12B3CF25U, 0x3182E51AU, 0x3C89EB13U, 0x2B94F908U, 0x269FF701U, - 0xBD464DE6U, 0xB04D43EFU, 0xA75051F4U, 0xAA5B5FFDU, 0x896A75C2U, 0x84617BCBU, 0x937C69D0U, 0x9E7767D9U, - 0xD51E3DAEU, 0xD81533A7U, 0xCF0821BCU, 0xC2032FB5U, 0xE132058AU, 0xEC390B83U, 0xFB241998U, 0xF62F1791U, - 0xD68D764DU, 0xDB867844U, 0xCC9B6A5FU, 0xC1906456U, 0xE2A14E69U, 0xEFAA4060U, 0xF8B7527BU, 0xF5BC5C72U, - 0xBED50605U, 0xB3DE080CU, 0xA4C31A17U, 0xA9C8141EU, 0x8AF93E21U, 0x87F23028U, 0x90EF2233U, 0x9DE42C3AU, - 0x063D96DDU, 0x0B3698D4U, 0x1C2B8ACFU, 0x112084C6U, 0x3211AEF9U, 0x3F1AA0F0U, 0x2807B2EBU, 0x250CBCE2U, - 0x6E65E695U, 0x636EE89CU, 0x7473FA87U, 0x7978F48EU, 0x5A49DEB1U, 0x5742D0B8U, 0x405FC2A3U, 0x4D54CCAAU, - 0xDAF741ECU, 0xD7FC4FE5U, 0xC0E15DFEU, 0xCDEA53F7U, 0xEEDB79C8U, 0xE3D077C1U, 0xF4CD65DAU, 0xF9C66BD3U, - 0xB2AF31A4U, 0xBFA43FADU, 0xA8B92DB6U, 0xA5B223BFU, 0x86830980U, 0x8B880789U, 0x9C951592U, 0x919E1B9BU, - 0x0A47A17CU, 0x074CAF75U, 0x1051BD6EU, 0x1D5AB367U, 0x3E6B9958U, 0x33609751U, 0x247D854AU, 0x29768B43U, - 0x621FD134U, 0x6F14DF3DU, 0x7809CD26U, 0x7502C32FU, 0x5633E910U, 0x5B38E719U, 0x4C25F502U, 0x412EFB0BU, - 0x618C9AD7U, 0x6C8794DEU, 0x7B9A86C5U, 0x769188CCU, 0x55A0A2F3U, 0x58ABACFAU, 0x4FB6BEE1U, 0x42BDB0E8U, - 0x09D4EA9FU, 0x04DFE496U, 0x13C2F68DU, 0x1EC9F884U, 0x3DF8D2BBU, 0x30F3DCB2U, 0x27EECEA9U, 0x2AE5C0A0U, - 0xB13C7A47U, 0xBC37744EU, 0xAB2A6655U, 0xA621685CU, 0x85104263U, 0x881B4C6AU, 0x9F065E71U, 0x920D5078U, - 0xD9640A0FU, 0xD46F0406U, 0xC372161DU, 0xCE791814U, 0xED48322BU, 0xE0433C22U, 0xF75E2E39U, 0xFA552030U, - 0xB701EC9AU, 0xBA0AE293U, 0xAD17F088U, 0xA01CFE81U, 0x832DD4BEU, 0x8E26DAB7U, 0x993BC8ACU, 0x9430C6A5U, - 0xDF599CD2U, 0xD25292DBU, 0xC54F80C0U, 0xC8448EC9U, 0xEB75A4F6U, 0xE67EAAFFU, 0xF163B8E4U, 0xFC68B6EDU, - 0x67B10C0AU, 0x6ABA0203U, 0x7DA71018U, 0x70AC1E11U, 0x539D342EU, 0x5E963A27U, 0x498B283CU, 0x44802635U, - 0x0FE97C42U, 0x02E2724BU, 0x15FF6050U, 0x18F46E59U, 0x3BC54466U, 0x36CE4A6FU, 0x21D35874U, 0x2CD8567DU, - 0x0C7A37A1U, 0x017139A8U, 0x166C2BB3U, 0x1B6725BAU, 0x38560F85U, 0x355D018CU, 0x22401397U, 0x2F4B1D9EU, - 0x642247E9U, 0x692949E0U, 0x7E345BFBU, 0x733F55F2U, 0x500E7FCDU, 0x5D0571C4U, 0x4A1863DFU, 0x47136DD6U, - 0xDCCAD731U, 0xD1C1D938U, 0xC6DCCB23U, 0xCBD7C52AU, 0xE8E6EF15U, 0xE5EDE11CU, 0xF2F0F307U, 0xFFFBFD0EU, - 0xB492A779U, 0xB999A970U, 0xAE84BB6BU, 0xA38FB562U, 0x80BE9F5DU, 0x8DB59154U, 0x9AA8834FU, 0x97A38D46U -}; - -__constant unsigned int USi3[256] = { - 0x00000000U, 0x090D0B0EU, 0x121A161CU, 0x1B171D12U, 0x24342C38U, 0x2D392736U, 0x362E3A24U, 0x3F23312AU, - 0x48685870U, 0x4165537EU, 0x5A724E6CU, 0x537F4562U, 0x6C5C7448U, 0x65517F46U, 0x7E466254U, 0x774B695AU, - 0x90D0B0E0U, 0x99DDBBEEU, 0x82CAA6FCU, 0x8BC7ADF2U, 0xB4E49CD8U, 0xBDE997D6U, 0xA6FE8AC4U, 0xAFF381CAU, - 0xD8B8E890U, 0xD1B5E39EU, 0xCAA2FE8CU, 0xC3AFF582U, 0xFC8CC4A8U, 0xF581CFA6U, 0xEE96D2B4U, 0xE79BD9BAU, - 0x3BBB7BDBU, 0x32B670D5U, 0x29A16DC7U, 0x20AC66C9U, 0x1F8F57E3U, 0x16825CEDU, 0x0D9541FFU, 0x04984AF1U, - 0x73D323ABU, 0x7ADE28A5U, 0x61C935B7U, 0x68C43EB9U, 0x57E70F93U, 0x5EEA049DU, 0x45FD198FU, 0x4CF01281U, - 0xAB6BCB3BU, 0xA266C035U, 0xB971DD27U, 0xB07CD629U, 0x8F5FE703U, 0x8652EC0DU, 0x9D45F11FU, 0x9448FA11U, - 0xE303934BU, 0xEA0E9845U, 0xF1198557U, 0xF8148E59U, 0xC737BF73U, 0xCE3AB47DU, 0xD52DA96FU, 0xDC20A261U, - 0x766DF6ADU, 0x7F60FDA3U, 0x6477E0B1U, 0x6D7AEBBFU, 0x5259DA95U, 0x5B54D19BU, 0x4043CC89U, 0x494EC787U, - 0x3E05AEDDU, 0x3708A5D3U, 0x2C1FB8C1U, 0x2512B3CFU, 0x1A3182E5U, 0x133C89EBU, 0x082B94F9U, 0x01269FF7U, - 0xE6BD464DU, 0xEFB04D43U, 0xF4A75051U, 0xFDAA5B5FU, 0xC2896A75U, 0xCB84617BU, 0xD0937C69U, 0xD99E7767U, - 0xAED51E3DU, 0xA7D81533U, 0xBCCF0821U, 0xB5C2032FU, 0x8AE13205U, 0x83EC390BU, 0x98FB2419U, 0x91F62F17U, - 0x4DD68D76U, 0x44DB8678U, 0x5FCC9B6AU, 0x56C19064U, 0x69E2A14EU, 0x60EFAA40U, 0x7BF8B752U, 0x72F5BC5CU, - 0x05BED506U, 0x0CB3DE08U, 0x17A4C31AU, 0x1EA9C814U, 0x218AF93EU, 0x2887F230U, 0x3390EF22U, 0x3A9DE42CU, - 0xDD063D96U, 0xD40B3698U, 0xCF1C2B8AU, 0xC6112084U, 0xF93211AEU, 0xF03F1AA0U, 0xEB2807B2U, 0xE2250CBCU, - 0x956E65E6U, 0x9C636EE8U, 0x877473FAU, 0x8E7978F4U, 0xB15A49DEU, 0xB85742D0U, 0xA3405FC2U, 0xAA4D54CCU, - 0xECDAF741U, 0xE5D7FC4FU, 0xFEC0E15DU, 0xF7CDEA53U, 0xC8EEDB79U, 0xC1E3D077U, 0xDAF4CD65U, 0xD3F9C66BU, - 0xA4B2AF31U, 0xADBFA43FU, 0xB6A8B92DU, 0xBFA5B223U, 0x80868309U, 0x898B8807U, 0x929C9515U, 0x9B919E1BU, - 0x7C0A47A1U, 0x75074CAFU, 0x6E1051BDU, 0x671D5AB3U, 0x583E6B99U, 0x51336097U, 0x4A247D85U, 0x4329768BU, - 0x34621FD1U, 0x3D6F14DFU, 0x267809CDU, 0x2F7502C3U, 0x105633E9U, 0x195B38E7U, 0x024C25F5U, 0x0B412EFBU, - 0xD7618C9AU, 0xDE6C8794U, 0xC57B9A86U, 0xCC769188U, 0xF355A0A2U, 0xFA58ABACU, 0xE14FB6BEU, 0xE842BDB0U, - 0x9F09D4EAU, 0x9604DFE4U, 0x8D13C2F6U, 0x841EC9F8U, 0xBB3DF8D2U, 0xB230F3DCU, 0xA927EECEU, 0xA02AE5C0U, - 0x47B13C7AU, 0x4EBC3774U, 0x55AB2A66U, 0x5CA62168U, 0x63851042U, 0x6A881B4CU, 0x719F065EU, 0x78920D50U, - 0x0FD9640AU, 0x06D46F04U, 0x1DC37216U, 0x14CE7918U, 0x2BED4832U, 0x22E0433CU, 0x39F75E2EU, 0x30FA5520U, - 0x9AB701ECU, 0x93BA0AE2U, 0x88AD17F0U, 0x81A01CFEU, 0xBE832DD4U, 0xB78E26DAU, 0xAC993BC8U, 0xA59430C6U, - 0xD2DF599CU, 0xDBD25292U, 0xC0C54F80U, 0xC9C8448EU, 0xF6EB75A4U, 0xFFE67EAAU, 0xE4F163B8U, 0xEDFC68B6U, - 0x0A67B10CU, 0x036ABA02U, 0x187DA710U, 0x1170AC1EU, 0x2E539D34U, 0x275E963AU, 0x3C498B28U, 0x35448026U, - 0x420FE97CU, 0x4B02E272U, 0x5015FF60U, 0x5918F46EU, 0x663BC544U, 0x6F36CE4AU, 0x7421D358U, 0x7D2CD856U, - 0xA10C7A37U, 0xA8017139U, 0xB3166C2BU, 0xBA1B6725U, 0x8538560FU, 0x8C355D01U, 0x97224013U, 0x9E2F4B1DU, - 0xE9642247U, 0xE0692949U, 0xFB7E345BU, 0xF2733F55U, 0xCD500E7FU, 0xC45D0571U, 0xDF4A1863U, 0xD647136DU, - 0x31DCCAD7U, 0x38D1C1D9U, 0x23C6DCCBU, 0x2ACBD7C5U, 0x15E8E6EFU, 0x1CE5EDE1U, 0x07F2F0F3U, 0x0EFFFBFDU, - 0x79B492A7U, 0x70B999A9U, 0x6BAE84BBU, 0x62A38FB5U, 0x5D80BE9FU, 0x548DB591U, 0x4F9AA883U, 0x4697A38DU -}; - -__constant unsigned int RCON[] = { - 0x01000000, 0x02000000, 0x04000000, 0x08000000, - 0x10000000, 0x20000000, 0x40000000, 0x80000000, - 0x1B000000, 0x36000000, // for 128-bit blocks, Rijndael never uses more than 10 rcon values -}; - -__constant unsigned char SBOXI[256] = -{ - 0x52, 0x09, 0x6A, 0xD5, 0x30, 0x36, 0xA5, 0x38, 0xBF, 0x40, 0xA3, 0x9E, 0x81, 0xF3, 0xD7, 0xFB, - 0x7C, 0xE3, 0x39, 0x82, 0x9B, 0x2F, 0xFF, 0x87, 0x34, 0x8E, 0x43, 0x44, 0xC4, 0xDE, 0xE9, 0xCB, - 0x54, 0x7B, 0x94, 0x32, 0xA6, 0xC2, 0x23, 0x3D, 0xEE, 0x4C, 0x95, 0x0B, 0x42, 0xFA, 0xC3, 0x4E, - 0x08, 0x2E, 0xA1, 0x66, 0x28, 0xD9, 0x24, 0xB2, 0x76, 0x5B, 0xA2, 0x49, 0x6D, 0x8B, 0xD1, 0x25, - 0x72, 0xF8, 0xF6, 0x64, 0x86, 0x68, 0x98, 0x16, 0xD4, 0xA4, 0x5C, 0xCC, 0x5D, 0x65, 0xB6, 0x92, - 0x6C, 0x70, 0x48, 0x50, 0xFD, 0xED, 0xB9, 0xDA, 0x5E, 0x15, 0x46, 0x57, 0xA7, 0x8D, 0x9D, 0x84, - 0x90, 0xD8, 0xAB, 0x00, 0x8C, 0xBC, 0xD3, 0x0A, 0xF7, 0xE4, 0x58, 0x05, 0xB8, 0xB3, 0x45, 0x06, - 0xD0, 0x2C, 0x1E, 0x8F, 0xCA, 0x3F, 0x0F, 0x02, 0xC1, 0xAF, 0xBD, 0x03, 0x01, 0x13, 0x8A, 0x6B, - 0x3A, 0x91, 0x11, 0x41, 0x4F, 0x67, 0xDC, 0xEA, 0x97, 0xF2, 0xCF, 0xCE, 0xF0, 0xB4, 0xE6, 0x73, - 0x96, 0xAC, 0x74, 0x22, 0xE7, 0xAD, 0x35, 0x85, 0xE2, 0xF9, 0x37, 0xE8, 0x1C, 0x75, 0xDF, 0x6E, - 0x47, 0xF1, 0x1A, 0x71, 0x1D, 0x29, 0xC5, 0x89, 0x6F, 0xB7, 0x62, 0x0E, 0xAA, 0x18, 0xBE, 0x1B, - 0xFC, 0x56, 0x3E, 0x4B, 0xC6, 0xD2, 0x79, 0x20, 0x9A, 0xDB, 0xC0, 0xFE, 0x78, 0xCD, 0x5A, 0xF4, - 0x1F, 0xDD, 0xA8, 0x33, 0x88, 0x07, 0xC7, 0x31, 0xB1, 0x12, 0x10, 0x59, 0x27, 0x80, 0xEC, 0x5F, - 0x60, 0x51, 0x7F, 0xA9, 0x19, 0xB5, 0x4A, 0x0D, 0x2D, 0xE5, 0x7A, 0x9F, 0x93, 0xC9, 0x9C, 0xEF, - 0xA0, 0xE0, 0x3B, 0x4D, 0xAE, 0x2A, 0xF5, 0xB0, 0xC8, 0xEB, 0xBB, 0x3C, 0x83, 0x53, 0x99, 0x61, - 0x17, 0x2B, 0x04, 0x7E, 0xBA, 0x77, 0xD6, 0x26, 0xE1, 0x69, 0x14, 0x63, 0x55, 0x21, 0x0C, 0x7D -}; -*/ - - #define MIN(a,b) (((a)<(b))?(a):(b)) #define AES_CTX_LENGTH 256 #define FALSE 0 @@ -478,7 +167,6 @@ __constant unsigned char SBOXI[256] = #define NONCE_SIZE 12 #define IV_SIZE 16 #define VMK_DECRYPT_SIZE 16 -#define VMK_SIZE 44 #define DICT_BUFSIZE (50*1024*1024) #define HASH_SIZE 8 //32 @@ -549,7 +237,7 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * unsigned int IV0, unsigned int IV4, unsigned int IV8, unsigned int IV12) { - int globalIndexPassword = (int)get_global_id(0); + int gIndex = (int)get_global_id(0); unsigned int hash0; unsigned int hash1; @@ -604,10 +292,10 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * unsigned int first_hash6; unsigned int first_hash7; - unsigned int indexW=(globalIndexPassword*FIXED_PASSWORD_BUFFER); + unsigned int indexW=(gIndex*FIXED_PASSWORD_BUFFER); int curr_fetch=0; - while(globalIndexPassword < numPassword) + while(gIndex < numPassword) { first_hash0 = 0x6A09E667; first_hash1 = 0xBB67AE85; @@ -628,13 +316,13 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * g = 0x1F83D9AB; h = 0x5BE0CD19; - indexW=(globalIndexPassword*FIXED_PASSWORD_BUFFER); + indexW=(gIndex*FIXED_PASSWORD_BUFFER); curr_fetch=0; index_generic=MAX_INPUT_PASSWORD_LEN; //--------------------- SCHEDULE ------------------- - //printf("MyId: %d, pass: %s, numPassword: %d\n", globalIndexPassword, w_password+indexW, numPassword); + //printf("MyId: %d, pass: %s, numPassword: %d\n", gIndex, w_password+indexW, numPassword); 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; @@ -1110,11 +798,9 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * indexW += SINGLE_BLOCK_W_SIZE; } -//BUG TITAN X! + //BUG TITAN X! for(index_generic=ITERATION_NUMBER/2; index_generic < ITERATION_NUMBER; index_generic++) { - - //Prima parte a = 0x6A09E667; b = 0xBB67AE85; c = 0x3C6EF372; @@ -1221,7 +907,6 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * hash6 = 0x1F83D9AB + g; hash7 = 0x5BE0CD19 + h; - //Seconda parte a = hash0; b = hash1; c = hash2; @@ -1231,7 +916,6 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * g = hash6; h = hash7; - //I primi 4 valori dei blocchi W sono sempre uguali 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) @@ -1310,19 +994,10 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * indexW += SINGLE_BLOCK_W_SIZE; } -//----------------------------------------------------- EXP KEY 256 ------------------------------------------------ - // with Nb=4 and Nk=256 -> Nr=15 so (14+1)*Nb=60 32b-words are needed - // 48 words - //----- 1 - // AddRoundKey - - /* REUSE OF SCHEDULE VARIABLES */ - schedule0 = IV0 ^ hash0; schedule1 = IV4 ^ hash1; schedule2 = IV8 ^ hash2; schedule3 = IV12 ^ 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); @@ -1352,7 +1027,6 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6); schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7); - //----- 2 hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^ (TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^ (TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^ @@ -1377,8 +1051,6 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6); schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7); - - //----- 3 hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^ (TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^ (TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^ @@ -1404,8 +1076,6 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6); schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7); - - //----- 4 hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^ (TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^ (TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^ @@ -1430,8 +1100,6 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6); schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7); - - //----- 5 hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^ (TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^ (TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^ @@ -1456,8 +1124,6 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6); schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7); - - //----- 6 hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^ (TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^ (TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^ @@ -1482,7 +1148,6 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6); schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7); - // last 4 words hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^ (TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^ (TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^ @@ -1491,7 +1156,6 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * hash2 ^= hash1; hash3 ^= hash2; - // NR-th round schedule0 = (TS2[(schedule4 >> 24) ] & 0xFF000000) ^ (TS3[(schedule5 >> 16) & 0xFF] & 0x00FF0000) ^ (TS0[(schedule6 >> 8) & 0xFF] & 0x0000FF00) ^ @@ -1512,24 +1176,27 @@ __kernel void opencl_bitcracker_attack(int numPassword, __global unsigned char * (TS0[(schedule5 >> 8) & 0xFF] & 0x0000FF00) ^ (TS1[(schedule6 ) & 0xFF] & 0x000000FF) ^ hash3; - schedule4 = (unsigned int )(((unsigned int )(schedule0 & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(schedule0 & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(schedule0 & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(schedule0 & 0x000000ff) << 24); schedule5 = (unsigned int )(((unsigned int )(schedule1 & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(schedule1 & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(schedule1 & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(schedule1 & 0x000000ff) << 24); schedule6 = (unsigned int )(((unsigned int )(schedule2 & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(schedule2 & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(schedule2 & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(schedule2 & 0x000000ff) << 24); schedule7 = (unsigned int )(((unsigned int )(schedule3 & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(schedule3 & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(schedule3 & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(schedule3 & 0x000000ff) << 24); if ( - ((vmkKey[0] ^ ((unsigned char) schedule4)) == VMK_SIZE) && + ((vmkKey[0] ^ ((unsigned char) schedule4)) == 0x2c) && ((vmkKey[1] ^ ((unsigned char) (schedule4 >> 8))) == 0x00) && - ((vmkKey[8] ^ ((unsigned char) schedule6)) <= 0x05) && - ((vmkKey[9] ^ ((unsigned char) (schedule6 >> 8))) == 0x20) - ) - { - found[0] = globalIndexPassword; - break; +#if STRICT_CHECK == 1 + ((vmkKey[8] ^ ((unsigned char) schedule6)) == 0x03) && +#else + ((vmkKey[8] ^ ((unsigned char) schedule6)) <= 0x05) && +#endif + ((vmkKey[9] ^ ((unsigned char) (schedule6 >> 8))) == 0x20) + ) + { + found[0] = gIndex; + break; } - globalIndexPassword += get_global_size(0); + gIndex += get_global_size(0); } return; diff --git a/OpenCL_version/main.c b/OpenCL_version/main.c index 9eb5eaf..0678e37 100755 --- a/OpenCL_version/main.c +++ b/OpenCL_version/main.c @@ -40,6 +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; // OpenCL Vars cl_context cxGPUContext; // OpenCL context @@ -51,14 +52,16 @@ cl_device_id* cdDevices; // OpenCL device(s) void usage(char *name) { - printf("\nUsage: %s -i -d \n\n" + printf("\nUsage: %s -f -d \n\n" "Options:\n\n" " -h, --help" "\t\t\tShow this help\n" - " -i, --diskimage" - "\t\tPath to your disk image\n" + " -f, --hashfile" + "\t\tPath to your input hash file (HashExtractor output)\n" " -d, --dictionary file" "\t\tPath to dictionary or alphabet file\n" + " -s, --strict" + "\t\tStrict check (use only in case of false positives)\n" " -p, --platform" "\t\tPlatform\n" " -g, --gpu" @@ -122,7 +125,7 @@ int checkDeviceStatistics() clGetDeviceInfo(devices[j], CL_DEVICE_VERSION, 0, NULL, &valueSize); value = (char*) malloc(valueSize); clGetDeviceInfo(devices[j], CL_DEVICE_VERSION, valueSize, value, NULL); - printf("Hardware version: %s\n", value); + printf("OpenCL version supported: %s\n", value); free(value); // print software driver version @@ -293,12 +296,8 @@ int destroyClCtx() int main (int argc, char **argv) { - char dictionaryFile[INPUT_SIZE]="\0"; - char diskImageFile[INPUT_SIZE]; - unsigned char * salt; - unsigned char * mac; - unsigned char * nonce; - unsigned char * encryptedVMK; + char * input_dictionary=NULL, * input_hash=NULL; + unsigned char *salt, *nonce, *vmk; uint32_t * w_blocks_d; long int totGlobalMem; @@ -332,7 +331,7 @@ int main (int argc, char **argv) {0, 0, 0, 0} }; - opt = getopt_long (argc, argv, "hi:d:t:b:p:g:", long_options, &option_index); + opt = getopt_long (argc, argv, "hf:d:t:b:p:g:s", long_options, &option_index); if (opt == -1) break; switch (opt) { @@ -340,13 +339,14 @@ int main (int argc, char **argv) usage(argv[0]); exit(EXIT_FAILURE); break; - case 'i': + case 'f': if(strlen(optarg) >= INPUT_SIZE) { - fprintf(stderr, "ERROR: Disk image path is bigger than %d\n", INPUT_SIZE); + fprintf(stderr, "ERROR: Inut hash file path is bigger than %d\n", INPUT_SIZE); exit(EXIT_FAILURE); } - strncpy(diskImageFile,optarg, strlen(optarg)+1); + input_hash=(char *)Calloc(INPUT_SIZE, sizeof(char)); + strncpy(input_hash, optarg, strlen(optarg)+1); break; case 'd': if(strlen(optarg) >= INPUT_SIZE) @@ -354,7 +354,8 @@ int main (int argc, char **argv) fprintf(stderr, "ERROR: Dictionary file path is bigger than %d\n", INPUT_SIZE); exit(EXIT_FAILURE); } - strncpy(dictionaryFile,optarg, strlen(optarg)+1); + input_dictionary=(char *)Calloc(INPUT_SIZE, sizeof(char)); + strncpy(input_dictionary,optarg, strlen(optarg)+1); break; case 't': psw_x_thread = atoi(optarg); @@ -373,6 +374,9 @@ int main (int argc, char **argv) case 'p': platform_id = atoi(optarg); break; + case 's': + strict_check = 1; + break; default: exit(EXIT_FAILURE); } @@ -386,14 +390,14 @@ int main (int argc, char **argv) exit(EXIT_FAILURE); } - if (dictionaryFile[0] == '\0'){ + if (input_dictionary == NULL){ printf("Missing dictionary file!\n"); usage(argv[0]); exit(EXIT_FAILURE); } - if (diskImageFile[0] == '\0'){ - printf("Missing dick image file!\n"); + if (input_hash == NULL){ + printf("Missing input hash file!\n"); usage(argv[0]); exit(EXIT_FAILURE); } @@ -413,13 +417,9 @@ int main (int argc, char **argv) //****************** Data from target file ******************* printf("\n====================================\nExtracting data from disk image\n====================================\n\n"); - salt = (unsigned char *) Calloc(SALT_SIZE, sizeof(unsigned char)); - mac = (unsigned char *) Calloc(MAC_SIZE, sizeof(unsigned char)); - nonce = (unsigned char *) Calloc(NONCE_SIZE, sizeof(unsigned char)); - encryptedVMK = (unsigned char *) Calloc(VMK_SIZE, sizeof(unsigned char)); - if(readData(diskImageFile, &salt, &mac, &nonce, &encryptedVMK) == BIT_FAILURE) + if(parse_data(input_hash, &salt, &nonce, &vmk) == BIT_FAILURE) { - fprintf(stderr, "Disk image error... exit!\n"); + fprintf(stderr, "Input hash format error... exit!\n"); goto cleanup; } //************************************************************ @@ -435,7 +435,7 @@ int main (int argc, char **argv) //********************************************** //************* Dictionary Attack ************* - opencl_attack(dictionaryFile, w_blocks_h, encryptedVMK, nonce, gridBlocks); + opencl_attack(input_dictionary, w_blocks_h, vmk, nonce, gridBlocks); //********************************************* cleanup: diff --git a/OpenCL_version/opencl_attack.c b/OpenCL_version/opencl_attack.c index 0bda148..6388b57 100755 --- a/OpenCL_version/opencl_attack.c +++ b/OpenCL_version/opencl_attack.c @@ -28,7 +28,7 @@ unsigned char outPsw[MAX_INPUT_PASSWORD_LEN+2]; int *hostFound, match; -unsigned char *hostPassword; +char *hostPassword; static int check_match() { int i=0; @@ -36,7 +36,7 @@ static int check_match() { if (hostFound[0] >= 0){ snprintf((char*)outPsw, MAX_INPUT_PASSWORD_LEN+1, "%s", hostPassword+(hostFound[0]*FIXED_PASSWORD_BUFFER) ); for(i=0; i= 5) - CC_SM50=1; CC_SM50=0; - if(DEV_NVIDIA == 1 && CC_SM50 == 1) - ciErr1 = clBuildProgram(cpProgram, 1, &(cdDevices[gpu_id]), "-I . -cl-nv-verbose -D DEV_NVIDIA_SM50=1", NULL, NULL); - else if(DEV_NVIDIA == 1) - ciErr1 = clBuildProgram(cpProgram, 1, &(cdDevices[gpu_id]), "-I . -cl-nv-verbose -D DEV_NVIDIA_SM50=0", NULL, NULL); - else - ciErr1 = clBuildProgram(cpProgram, 1, &(cdDevices[gpu_id]), "-I . -D DEV_NVIDIA_SM50=0", NULL, NULL); + if(ccMajor >= 5) CC_SM50=1; + memset(optProgram, 0, 128); + if(DEV_NVIDIA == 1) + snprintf(optProgram, 128, "-I . -cl-nv-verbose -D DEV_NVIDIA_SM50=%d -D STRICT_CHECK=%d", CC_SM50, strict_check); + else + snprintf(optProgram, 128, "-I . -D DEV_NVIDIA_SM50=0 -D STRICT_CHECK=%d", strict_check); + + ciErr1 = clBuildProgram(cpProgram, 1, &(cdDevices[gpu_id]), optProgram, NULL, NULL); CL_ERROR(ciErr1); ret_cl = clGetProgramBuildInfo(cpProgram, cdDevices[gpu_id], CL_PROGRAM_BUILD_LOG, 0, NULL, &len); @@ -177,7 +175,7 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp numPassword = GPU_MAX_WORKGROUP_SIZE*gridBlocks*MAX_PASSWD_SINGLE_KERNEL; passwordBufferSize = numPassword * MAX_INPUT_PASSWORD_LEN * sizeof(unsigned char); - hostPassword = (unsigned char *) calloc(passwordBufferSize, sizeof(unsigned char)); + hostPassword = (char *) calloc(passwordBufferSize, sizeof(char)); hostFound = (int *) calloc(1, sizeof(int)); // -------------------------------------------------------------------------- @@ -187,9 +185,6 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp deviceEncryptedVMK = clCreateBuffer(cxGPUContext, CL_MEM_READ_WRITE, VMK_DECRYPT_SIZE*sizeof(unsigned char), NULL, &ciErr1); CL_ERROR(ciErr1); - // deviceIV = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, IV_SIZE*sizeof(unsigned char), NULL, &ciErr1); - //CL_ERROR(ciErr1); - devicePassword = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, passwordBufferSize*sizeof(unsigned char), NULL, &ciErr1); CL_ERROR(ciErr1); @@ -200,16 +195,7 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp CL_ERROR(ciErr1); // -------------------------------------------------------------------------- - - // ------------------------------- Write static buffers ------------------------------- - /* - schedule0 = __byte_perm(((unsigned int *)(IV))[0], 0, 0x0123) ^ hash0; - schedule1 = __byte_perm(((unsigned int *)(IV+4))[0], 0, 0x0123) ^ hash1; - schedule2 = __byte_perm(((unsigned int *)(IV+8))[0], 0, 0x0123) ^ hash2; - schedule3 = __byte_perm(((unsigned int *)(IV+12))[0], 0, 0x0123) ^ hash3; - */ - 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); @@ -225,9 +211,8 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp 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, 0, NULL, NULL); - CL_ERROR(ciErr1); - + ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, deviceEncryptedVMK, CL_TRUE, 0, VMK_DECRYPT_SIZE*sizeof(char), encryptedVMK+16, 0, NULL, NULL); + CL_ERROR(ciErr1); // -------------------------------------------------------------------------- szLocalWorkSize = GPU_MAX_WORKGROUP_SIZE; @@ -239,7 +224,6 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp int iter=0; while(!feof(fp_file_passwords)) { - numReadPassword = readFilePassword(&hostPassword, numPassword, fp_file_passwords); /* Copy input data to memory buffer */ @@ -278,9 +262,7 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp ciErr1 |= clSetKernelArg(ckKernelAttack, 8, sizeof(cl_int), (void*)&IV12); CL_ERROR(ciErr1); // -------------------------------------------------------- - - // Launch kernel - + time_t start,end; double dif; TIMER_DEF(0); diff --git a/OpenCL_version/utils.c b/OpenCL_version/utils.c index 76292f4..42c0868 100755 --- a/OpenCL_version/utils.c +++ b/OpenCL_version/utils.c @@ -21,6 +21,25 @@ #include "bitcracker.h" +/* John The Ripper function */ +char *strtokm(char *s1, const char *delims) +{ + static char *last = NULL; + char *endp; + + if (!s1) + s1 = last; + if (!s1 || *s1 == 0) + return last = NULL; + endp = strpbrk(s1, delims); + if (endp) { + *endp = '\0'; + last = endp + 1; + } else + last = NULL; + return s1; +} + void * Calloc(size_t len, size_t size) { void * ptr = NULL; if( size <= 0) @@ -54,110 +73,116 @@ void print_hex(unsigned char *str, int len) printf("%02x", str[i]); } -int readData(char * encryptedImagePath, unsigned char ** salt, unsigned char ** mac, unsigned char ** nonce, unsigned char ** encryptedVMK) +int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, unsigned char ** vmk) { - int match = 0; - const char signature[9] = "-FVE-FS-"; - int version = 0; - unsigned char vmk_entry[4] = { 0x02, 0x00, 0x08, 0x00 }; - unsigned char key_protection_type[2] = { 0x00, 0x20 }; - unsigned char value_type[2] = { 0x00, 0x05 }; - char c; - int i = 0; - int j, fileLen; + char * hash; + char *p; + int i, salt_len, iterations, vmk_size, nonce_len; + FILE * fphash; + char tmp[2]; + int j=0; - if( !salt || !mac || !nonce || !encryptedVMK ) { - fprintf(stderr, "Input error\n"); - return BIT_FAILURE; + (*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)); + hash = (char *) Calloc(INPUT_HASH_SIZE, sizeof(char)); + + if(!input_hash) + { + fprintf(stderr, "No input hash provided\n"); + goto out; } - printf("Opening file %s\n", encryptedImagePath); - FILE * encryptedImage = fopen(encryptedImagePath, "r"); - if (!encryptedImage) { - fprintf(stderr, "! %s : %s\n", encryptedImagePath, strerror(errno)); - return BIT_FAILURE; + fphash = fopen(input_hash, "r"); + if (!fphash) { + fprintf(stderr, "! %s : %s\n", input_hash, strerror(errno)); + goto out; + } + fgets(hash, INPUT_HASH_SIZE, fphash); + fclose(fphash); + if(!hash) + { + fprintf(stderr, "No correct input hash provided\n"); + goto out; } - fseek(encryptedImage, 0, SEEK_END); - fileLen = ftell(encryptedImage); - fseek(encryptedImage, 0, SEEK_SET); - for (j = 0; j < fileLen; j++) { - c = fgetc(encryptedImage); - while (i < 8 && (unsigned char)c == signature[i]) { - c = fgetc(encryptedImage); - i++; - } - if (i == 8) { - match = 1; - fprintf(stderr, "Signature found at 0x%08lx\n", (ftell(encryptedImage) - i - 1)); - fseek(encryptedImage, 1, SEEK_CUR); - version = fgetc(encryptedImage); - fprintf(stderr, "Version: %d ", version); - if (version == 1) - fprintf(stderr, "(Windows Vista)\n"); - else if (version == 2) - fprintf(stderr, "(Windows 7 or later)\n"); - else { - fprintf(stderr, "\nInvalid version, looking for a signature with valid version...\n"); - } - } - i = 0; - while (i < 4 && (unsigned char)c == vmk_entry[i]) { - c = fgetc(encryptedImage); - i++; - } + printf("Hash file %s:\n%s\n", input_hash, hash); - if (i == 4) { - fprintf(stderr, "VMK entry found at 0x%08lx\n", (ftell(encryptedImage) - i - 3)); - fseek(encryptedImage, 27, SEEK_CUR); - if ( - ( (unsigned char)fgetc(encryptedImage) == key_protection_type[0]) && - ( (unsigned char)fgetc(encryptedImage) == key_protection_type[1]) - ) - { - fprintf(stderr, "Key protector with user password found\n"); - fseek(encryptedImage, 12, SEEK_CUR); - fillBuffer(encryptedImage, *salt, SALT_SIZE); - fseek(encryptedImage, 83, SEEK_CUR); - if (((unsigned char)fgetc(encryptedImage) != value_type[0]) || ((unsigned char)fgetc(encryptedImage) != value_type[1])) { - fprintf(stderr, "Error: VMK not encrypted with AES-CCM\n"); - //ret failure? - } - - fseek(encryptedImage, 3, SEEK_CUR); - fillBuffer(encryptedImage, *nonce, NONCE_SIZE); - fillBuffer(encryptedImage, *mac, MAC_SIZE); - fillBuffer(encryptedImage, *encryptedVMK, VMK_SIZE); - - fprintf(stdout, "Nonce:\n"); - print_hex(*nonce, NONCE_SIZE); - fprintf(stdout, "\n"); - - fprintf(stdout, "MAC:\n"); - print_hex(*mac, MAC_SIZE); - fprintf(stdout, "\n"); - - fprintf(stdout, "VMK:\n"); - print_hex(*encryptedVMK, VMK_SIZE); - fprintf(stdout, "\n"); - break; - } - } - - i = 0; + if (strncmp(hash, HASH_TAG, HASH_TAG_LEN) != 0) + { + fprintf(stderr, "Wrong hash format\n"); + goto out; } - fclose(encryptedImage); + hash += HASH_TAG_LEN; + + p = strtokm(hash, "$"); // version + p = strtokm(NULL, "$"); // salt length - if (match == 0) { - fprintf(stderr, "Error while extracting data: No signature found!\n"); - return BIT_FAILURE; + salt_len = atoi(p); + if(salt_len != SALT_SIZE) + { + fprintf(stderr, "Wrong Salt size\n"); + goto out; } + p = strtokm(NULL, "$"); // salt + for (i = 0, j = 0; i < salt_len*2; i+=2, j++) + { + tmp[0] = p[i]; + tmp[1] = p[i+1]; + long int ret = strtol(tmp, NULL, 16); + (*salt)[j] = (unsigned char)(ret); //((ARCH_INDEX(p[i * 2]) * 16) + ARCH_INDEX(p[i * 2 + 1])); + } + + p = strtokm(NULL, "$"); // iterations + iterations = atoi(p); + p = strtokm(NULL, "$"); // nonce length + nonce_len = atoi(p); + if(nonce_len != NONCE_SIZE) + { + fprintf(stderr, "Wrong Nonce size\n"); + goto out; + } + + p = strtokm(NULL, "$"); // nonce + for (i = 0, j = 0; i < NONCE_SIZE*2; i+=2, j++) + { + tmp[0] = p[i]; + tmp[1] = p[i+1]; + long int ret = strtol(tmp, NULL, 16); + (*nonce)[j] = (unsigned char)(ret); //((ARCH_INDEX(p[i * 2]) * 16) + ARCH_INDEX(p[i * 2 + 1])); + } + + p = strtokm(NULL, "$"); // data_size + + vmk_size = atoi(p); + if(vmk_size != VMK_SIZE) + { + fprintf(stderr, "Wrong VMK size\n"); + goto out; + } + + p = strtokm(NULL, "$"); // data + for (i = 0, j = 0; i < vmk_size*2; i+=2, j++) + { + tmp[0] = p[i]; + tmp[1] = p[i+1]; + long int ret = strtol(tmp, NULL, 16); + (*vmk)[j] = (unsigned char)(ret); //((ARCH_INDEX(p[i * 2]) * 16) + ARCH_INDEX(p[i * 2 + 1])); + } + return BIT_SUCCESS; + + out: + free(*salt); + free(*nonce); + free(*vmk); + + return BIT_FAILURE; } -int readFilePassword(unsigned char ** buf, int maxNumPsw, FILE *fp) { +int readFilePassword(char ** buf, int maxNumPsw, FILE *fp) { int i=0, size; char tmp[FIXED_PASSWORD_BUFFER]; memset(tmp, 0, FIXED_PASSWORD_BUFFER); diff --git a/README.md b/README.md index 8678954..5b56f30 100644 --- a/README.md +++ b/README.md @@ -1,7 +1,7 @@ -Bitcracker +BitCracker ======== -BitCracker is the first open source BitLocker password cracker tool +BitCracker is the first open source BitLocker password cracking tool Introduction === @@ -12,7 +12,7 @@ BitCracker is a mono-GPU password cracking tool for memory units encrypted with ![alt text](http://openwall.info/wiki/_media/john/bitcracker_img1.png) Our attack has been tested on several memory units encrypted with BitLocker running on Windows 7, Window 8.1 and Windows 10 (both compatible and non-compatible mode). -here we present two implementations: CUDA and OpenCL. +Here we present two implementations: CUDA and OpenCL. Requirements === @@ -22,33 +22,60 @@ For CUDA implementation, you need at least CUDA 7.5 and an NVIDIA GPU with minim How To === -Use the buil.sh script to run makefiles within CUDA_version and OpenCL_version. +Use the build.sh script to build 3 executables: + +- hash extractor +- BitCracker CUDA version +- BitCracker OpenCL version + +The executables are stored in the build directory. +
+Before starting the attack, you need to run bitcracker_hash to extract the hash from the encrypted memory unit. ``` -> ./bitcracker_cuda -h +> ./build/bitcracker_hash -h -Usage: ./bitcracker_cuda -i -d +Usage: ./build/bitcracker_hash -i -o Options: - -h, --help Show this help - -i, --diskimage Path to your disk image - -d, --dictionary file Path to dictionary or alphabet file - -g, --gpu GPU device number - -t, --passthread Set the number of password per thread threads - -b, --blocks Set the number of blocks + -h, --help Show this help + -i, --disk Path of memory unit encrypted with BitLocker + -o, --outfile Output file ``` -You can use the run.sh script to execute two simple runs using the encrypted images we provide in this repo: +The extracted hash is fully compatible with the John The Ripper format (see next Section).
+ +Then you can use the output hash file to run the BitCracker attack. + +``` +> ./build/bitcracker_cuda -h + +Usage: ./build/bitcracker_cuda -f -d + +Options: + + -h, --help Show this help + -f, --hashfile Path to your input hash file (HashExtractor output) + -s, --strict Strict check (use only in case of false positives) + -d, --dictionary Path to dictionary or alphabet file + -g, --gpu GPU device number + -t, --passthread Set the number of password per thread threads + -b, --blocks Set the number of blocks +``` + +N.B. In case of false positives, you can use the -s option (empirically improved password check). + +In the the run_test.sh script there are several attack examples using the encrypted images provided in this repo: * imgWin7: memory unit encrypted with BitLocker using Windows 7 Enteprise edition OS * imgWin8: memory unit encrypted with BitLocker using Windows 8 Enteprise edition OS * imgWin10Compatible.vhd: memory unit encrypted with BitLocker (compatible mode) using Windows 10 Enteprise edition OS, * imgWin10NonCompatible.vhd: memory unit encrypted with BitLocker (NON compatible mode) using Windows 10 Enteprise edition OS, -* imgWin10CompatibleLong27.vhd: memory unit encrypted with BitLocker (compatible mode) using Windows 10 Enteprise edition OS using the longest possible password +* imgWin10CompatibleLong27.vhd: memory unit encrypted with BitLocker (compatible mode) using Windows 10 Enteprise edition OS using the longest possible password (27 characters) -Currently, BitCracker is able to evaluate passwords having length between 8 (minimum password length) and 27 characters (implementation reasons). -In the next release, the maximum password lenght will be dynamic. -BitCracker doesn't provide any mask attack, cache mechanism or smart dictionary creation; thus you need to provide your own input dictionary. +Currently, BitCracker is able to evaluate passwords having length between 8 (minimum password length) and 27 characters (implementation reasons). + +BitCracker doesn't provide any mask attack, cache mechanism or smart dictionary creation; therefore you need to provide your own input dictionary. Performance === @@ -81,6 +108,13 @@ We released the OpenCL version as a plugin of the John The Ripper (bleeding jumb * Wiki page: http://openwall.info/wiki/john/OpenCL-BitLocker
* JtR source code: https://github.com/magnumripper/JohnTheRipper +Future Works +=== + +In the next relese: +- The maximum password lenght will be dynamic +- Improve strict check with optional MAC verification to avoid any false positive + References, credits and contacts === diff --git a/build.sh b/build.sh index 5d8e649..8c1cfbb 100755 --- a/build.sh +++ b/build.sh @@ -1,3 +1,20 @@ #!/usr/bin/env bash +mkdir -p ./build + +printf "\n====== Build BitCracker Hash Extractor ======\n" +cd Hash_Extractor && make clean && make +cd .. +mv Hash_Extractor/bitcracker_hash build 2> /dev/null + +printf "\n====== Build BitCracker CUDA version ======\n" cd CUDA_version && make clean && make +cd .. +mv CUDA_version/bitcracker_cuda build 2> /dev/null + +printf "\n====== Build BitCracker OpenCL version ======\n" +cd OpenCL_version && make clean && make +cd .. +mv OpenCL_version/bitcracker_opencl build 2> /dev/null + +printf "\n====== Executables in build directory ======\n" diff --git a/run.sh b/run.sh deleted file mode 100755 index fb5e262..0000000 --- a/run.sh +++ /dev/null @@ -1,19 +0,0 @@ -#!/usr/bin/env bash - -#Show help -./CUDA_version/bitcracker_cuda -h - -#Image encrypted with Windows 8.1 Enterprise -./CUDA_version/bitcracker_cuda -i ./Images/imgWin8 -d ./Dictionary/test_passwords.txt -t 1 -b 1 -g 0 - -#Image encrypted with Windows 7 Pro -./CUDA_version/bitcracker_cuda -i ./Images/imgWin7 -d ./Dictionary/test_passwords.txt -t 1 -b 1 -g 0 - -#Image encrypted with Windows 10 Enteprise using BitLocker compatible mode -./CUDA_version/bitcracker_cuda -i ./Images/imgWin10Compatible.vhd -d ./Dictionary/test_passwords.txt -t 1 -b 1 -g 0 - -#Image encrypted with Windows 10 Enteprise using BitLocker non-compatible mode -./CUDA_version/bitcracker_cuda -i ./Images/imgWin10NonCompatible.vhd -d ./Dictionary/test_passwords.txt -t 1 -b 1 -g 0 - -#Image encrypted with Windows 10 Enteprise using BitLocker non-compatible mode, long password (27 chars) -./CUDA_version/bitcracker_cuda -i ./Images/imgWin10NonCompatibleLong27.vhd -d ./Dictionary/test_passwords.txt -t 1 -b 1 -g 0 diff --git a/run_tests.sh b/run_tests.sh new file mode 100755 index 0000000..e230b8b --- /dev/null +++ b/run_tests.sh @@ -0,0 +1,40 @@ +#!/usr/bin/env bash + + +mkdir -p test_hash + +printf "\n\n************ Extract BitLocker hash from encrypted memory units ************\n\n" + +#Image encrypted with Windows 8.1 Enterprise +./build/bitcracker_hash -o test_hash/imgWin8.txt -i ./Images/imgWin8 +printf "\n\n" + +#Image encrypted with Windows 7 Pro +./build/bitcracker_hash -o test_hash/imgWin7.txt -i ./Images/imgWin7 +printf "\n\n" + +#Image encrypted with Windows 10 Enteprise using BitLocker compatible mode +./build/bitcracker_hash -o test_hash/imgWin10Compatible.txt -i ./Images/imgWin10Compatible.vhd +printf "\n\n" + +#Image encrypted with Windows 10 Enteprise using BitLocker non-compatible mode +./build/bitcracker_hash -o test_hash/imgWin10NonCompatible.txt -i ./Images/imgWin10NonCompatible.vhd +printf "\n\n" + +#Image encrypted with Windows 10 Enteprise using BitLocker non-compatible mode, long password (27 chars) +./build/bitcracker_hash -o test_hash/imgWin10NonCompatibleLong27.txt -i ./Images/imgWin10NonCompatibleLong27.vhd + + +printf "\n\n************ Testing BitCracker CUDA version ************\n\n" +#Show help +./build/bitcracker_cuda -h + +./build/bitcracker_cuda -f ./test_hash/imgWin8.txt -d ./Dictionary/test_passwords.txt -t 1 -b 1 -g 0 + +./build/bitcracker_cuda -f ./test_hash/imgWin7.txt -d ./Dictionary/test_passwords.txt -t 1 -b 1 -g 0 + +./build/bitcracker_cuda -f ./test_hash/imgWin10Compatible.txt -d ./Dictionary/test_passwords.txt -t 1 -b 1 -g 0 + +./build/bitcracker_cuda -f ./test_hash/imgWin10NonCompatible.txt -d ./Dictionary/test_passwords.txt -t 1 -b 1 -g 0 + +./build/bitcracker_cuda -f ./test_hash/imgWin10NonCompatibleLong27.txt -d ./Dictionary/test_passwords.txt -t 1 -b 1 -g 0