mirrors/bitcracker
2
0
Fork 0
mirror of https://github.com/e-ago/bitcracker.git synced 2025-10-27 07:29:18 +00:00
Code Issues Packages Projects Releases Wiki Activity

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).

Browse Source
This commit is contained in:
elenago 2017-09-19 10:59:03 +02:00
parent 7be094d24b
commit e9c8f26e2d
17 changed files with 767 additions and 730 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
CUDA_version/bitcracker.h
View File

@ -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);

93
CUDA_version/cuda_attack.cu
View File

@ -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<<<gridBlocks, ATTACK_DEFAULT_THREADS, 0, stream[indexStream]>>>(indexStream, numReadPassword[indexStream], deviceFound[indexStream], deviceEncryptedVMK, deviceIV);
decrypt_vmk<<<gridBlocks, ATTACK_DEFAULT_THREADS, 0, stream[indexStream]>>>(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;

58
CUDA_version/main.cu
View File

@ -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 <disk_image> -d <dictionary_file>\n\n"
printf("\nUsage: %s -f <hash_file> -d <dictionary_file>\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;
}

202
CUDA_version/utils.cu
View File

@ -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];

6
Hash_Extractor/Makefile Executable file
View File

@ -0,0 +1,6 @@
bitcracker_hash:
gcc bitcracker_hash.c -o bitcracker_hash
clean:
rm -rf *.o
rm -rf bitcracker_hash

BIN
Hash_Extractor/bitcracker_hash Executable file
View File

Binary file not shown.

290
Hash_Extractor/bitcracker_hash.c Normal file
View File

@ -0,0 +1,290 @@
/*
* BitCracker: BitLocker password cracking tool, Hash Extractor.
* Copyright (C) 2013-2017 Elena Ago <elena dot ago at gmail dot com>
* Massimo Bernaschi <massimo dot bernaschi at gmail dot com>
*
* 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 <http://www.gnu.org/licenses/>.
*/
#include <stdio.h>
#include <stdint.h>
#include <string.h>
#include <unistd.h>
#include <ctype.h>
#include <getopt.h>
#include <errno.h>
#include <limits.h>
#include <stdlib.h>
#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 <Encrypted memory unit> -o <output file>\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;
}

1
Hash_Extractor/imgWin10Compatible.txt Normal file
View File

@ -0,0 +1 @@
$bitlocker$0$16$91a4ec232ab95bbb8e8ef964308c6b47$1048576$12$306af9dca50fd30103000000$60$00000000000000000000000000000000509aab04f2161082ed6153d6ea8ad51d45c1ae6ae77cdc470789472640f409a1c2ede715ea5a6bbc320e2312

17
OpenCL_version/bitcracker.h
View File

@ -1,5 +1,5 @@
/*
* BitCracker: BitLocker password cracking tool, CUDA version.
* BitCracker: BitLocker password cracking tool, OpenCL version.
* Copyright (C) 2013-2017 Elena Ago <elena dot ago at gmail dot com>
* Massimo Bernaschi <massimo dot bernaschi at gmail dot com>
*
@ -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);

371
OpenCL_version/kernel_attack.cl
View File

@ -1,7 +1,7 @@
/*
* BitCracker: BitLocker password cracking tool, OpenCL version.
* Copyright (C) 2013-2017 Elena Ago <elena dot ago at gmail dot com>
* Massimo Bernaschi <massimo dot bernaschi at gmail dot com>
* Massimo Bernaschi <massimo dot bernaschi at gmail dot com>
*
* 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;

50
OpenCL_version/main.c
View File

@ -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 <disk_image> -d <dictionary_file>\n\n"
printf("\nUsage: %s -f <hash_file> -d <dictionary_file>\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:

48
OpenCL_version/opencl_attack.c
View File

@ -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<MAX_INPUT_PASSWORD_LEN; i++)
if(outPsw[i] == 0x80 || outPsw[i] == 0xffffff80) outPsw[i]='\0';
if(outPsw[i] == 0x80 || outPsw[i] == 0xff) outPsw[i]='\0'; //0xffffff80
return 1;
}
@ -63,6 +63,7 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp
char *source_str_attack;
size_t len = 0;
cl_int ret_cl = CL_SUCCESS, ret_cl_log = CL_SUCCESS, ret_info_kernel = CL_SUCCESS;
char optProgram[128];
//------- READ CL FILE ------
/* Load kernel source file */
@ -102,8 +103,6 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp
tmpIV[IV_SIZE-1] = 1;
// --------------------------------------------------------------------------
// ---- Open File Dictionary ----
if (!memcmp(dname, "-\0", 2)) {
fp_file_passwords= stdin;
@ -115,7 +114,6 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp
}
}
// --------------------------------------------------------------------------
// ------------------------------- Kernel setup -------------------------------
/* Create kernel from source */
@ -123,16 +121,16 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp
CL_ERROR(ciErr1);
clGetDeviceInfo(cdDevices[gpu_id], CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, sizeof(ccMajor), &ccMajor, NULL);
if(ccMajor >= 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);

203
OpenCL_version/utils.c
View File

@ -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);

68
README.md
View File

@ -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.
<br>
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 <disk_image> -d <dictionary_file>
Usage: ./build/bitcracker_hash -i <Encrypted memory unit> -o <output file>
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).<br>
Then you can use the output hash file to run the BitCracker attack.
```
> ./build/bitcracker_cuda -h
Usage: ./build/bitcracker_cuda -f <hash_file> -d <dictionary_file>
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 <br />
* 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
===

17
build.sh
View File

@ -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"

19
run.sh
View File

@ -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

40
run_tests.sh Executable file
View File

@ -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