Minor fixes, MAC option added (CUDA version only)

This commit is contained in:
elenago 2017-11-14 12:00:39 +01:00
parent 22e77eeca8
commit b28c8e6fdc
16 changed files with 1366 additions and 392 deletions

View File

@ -3,4 +3,4 @@ bitcracker_cuda:
clean:
rm -rf *.o
rm -rf bitcracker_cuda
rm -rf ../build/bitcracker_cuda

View File

@ -3,7 +3,7 @@
* 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.
* This file is part of the BitCracker project: https://github.com/e-ago/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

View File

@ -3,7 +3,7 @@
* 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.
* This file is part of the BitCracker project: https://github.com/e-ago/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
@ -27,6 +27,7 @@
#include <errno.h>
#include <limits.h>
#include <stdlib.h>
#include <string.h>
#include "sha256_header.h"
#include "aes_header.h"
@ -40,13 +41,16 @@
#define NONCE_SIZE 12
#define IV_SIZE 16
#define VMK_SIZE 60
#define VMK_DECRYPT_SIZE 16
#define VMK_HEADER_SIZE 12
#define VMK_BODY_SIZE 32
#define VMK_FULL_SIZE 44
#define DICT_BUFSIZE (50*1024*1024)
#define MAX_PLEN 32
#define HASH_TAG "$bitlocker$"
#define HASH_TAG_LEN (sizeof(HASH_TAG) - 1)
#define INPUT_HASH_SIZE 210
#define INPUT_HASH_SIZE 245
#ifndef UINT32_C
#define UINT32_C(c) c ## UL
#endif
@ -66,7 +70,8 @@
#define BLOCK_UNIT 32
#define HASH_SIZE_STRING 32
#define ATTACK_DEFAULT_THREADS 1024
#define CUDA_THREADS_NO_MAC 1024
#define CUDA_THREADS_WITH_MAC 256
#define BIT_SUCCESS 0
#define BIT_FAILURE 1
@ -91,17 +96,35 @@ extern int psw_x_thread;
extern int tot_psw;
extern size_t size_psw;
extern int strict_check;
extern int mac_comparison;
extern unsigned char * salt;
/* ++++++++++++++++++++++++++++++++++++++ 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, int check);
__global__ __launch_bounds__(1024,1) void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsigned char * vmkKey,
unsigned char * IV, int strict_check,
int v0, int v1, int v2, int v3,
uint32_t s0, uint32_t s1, uint32_t s2, uint32_t s3);
__global__ __launch_bounds__(1024,1) void decrypt_vmk_with_mac(
int numStream, int tot_psw_kernel, int *found,
unsigned char * vmkKey, unsigned char * vmkIV,
unsigned char * mac, unsigned char * macIV, unsigned char * computeMacIV,
int v0, int v1, int v2, int v3,
uint32_t s0, uint32_t s1, uint32_t s2, uint32_t s3
);
__global__ void decrypt_vmk_noimprovements(
int numStream, int tot_psw_kernel, int *found,
unsigned char * vmkKey, unsigned char * vmkIV,
unsigned char * mac, unsigned char * macIV, unsigned char * computeMacIV,
int v0, int v1, int v2, int v3,
uint32_t s0, uint32_t s1, uint32_t s2, uint32_t s3
);
/* ++++++++++++++++++++++++++++++++++++++ HOST FUNCTIONS ++++++++++++++++++++++++++++++++++++++ */
int w_block_precomputed(unsigned char * salt, uint32_t * w_blocks_d);
int readFilePassword(char ** buf, int maxNumPsw, FILE *fp);
int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, unsigned char ** vmk);
int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, unsigned char ** vmk, unsigned char ** mac);
char *strtokm(char *s1, const char *delims);
char *cuda_attack(char *dname, uint32_t * w_blocks_d, unsigned char * encryptedVMK, unsigned char * nonce, int gridBlocks);
char *cuda_attack(char *dname, uint32_t * w_blocks_d, unsigned char * encryptedVMK, unsigned char * nonce, unsigned char * encryptedMAC, int gridBlocks);
void setBufferPasswordSize(size_t avail, size_t * passwordBufferSize, int * numPassword);
void * Calloc(size_t len, size_t size);

File diff suppressed because it is too large Load Diff

View File

@ -3,7 +3,7 @@
* 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.
* This file is part of the BitCracker project: https://github.com/e-ago/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
@ -27,6 +27,8 @@ int tot_psw=0;
size_t size_psw=0;
size_t tot_word_mem=(SINGLE_BLOCK_SHA_SIZE * ITERATION_NUMBER * sizeof(uint32_t));
int strict_check=0;
int mac_comparison=0;
unsigned char * salt;
void usage(char *name){
printf("\nUsage: %s -f <hash_file> -d <dictionary_file>\n\n"
@ -67,8 +69,8 @@ int getGPUStats()
printf("Clock rate: %d\n", prop.clockRate );
printf("Clock rate: %.0f MHz (%.02f GHz)\n", prop.clockRate * 1e-3f, prop.clockRate * 1e-6f);
printf("Memory Clock Rate (KHz): %d\n", prop.memoryClockRate);
printf("Memory Bus Width (bits): %d\n", prop.memoryBusWidth);
printf("Peak Memory Bandwidth (GB/s): %f\n", 2.0*prop.memoryClockRate*(prop.memoryBusWidth/8)/1.0e6);
printf("Memory Bus Width (bits): %d\n", prop.memoryBusWidth);
printf("Peak Memory Bandwidth (GB/s): %f\n", 2.0*prop.memoryClockRate*(prop.memoryBusWidth/8)/1.0e6);
printf("Device copy overlap: " );
if (prop.deviceOverlap) printf("Enabled\n" ); else printf("Disabled\n" );
printf("Async memory engine count: %d\n", prop.asyncEngineCount);
@ -101,7 +103,7 @@ int getGPUStats()
int main (int argc, char **argv)
{
char * input_dictionary=NULL, * input_hash=NULL;
unsigned char *salt, *nonce, *vmk;
unsigned char *nonce, *vmk, *mac;
uint32_t * w_blocks_d;
int gridBlocks = 1;
@ -129,22 +131,12 @@ int main (int argc, char **argv)
{0, 0, 0, 0}
};
opt = getopt_long(argc, argv, "hf:d:t:b:g:s", long_options, &option_index);
opt = getopt_long(argc, argv, "b:d:f:g:mhst:", long_options, &option_index);
if (opt == -1)
break;
switch (opt) {
case 'h':
usage(argv[0]);
exit(EXIT_FAILURE);
break;
case 'f':
if(strlen(optarg) >= INPUT_SIZE)
{
fprintf(stderr, "ERROR: Inut hash file path is bigger than %d\n", INPUT_SIZE);
exit(EXIT_FAILURE);
}
input_hash=(char *)Calloc(INPUT_SIZE, sizeof(char));
strncpy(input_hash, optarg, strlen(optarg)+1);
case 'b':
gridBlocks = atoi(optarg);
break;
case 'd':
if(strlen(optarg) >= INPUT_SIZE)
@ -155,6 +147,28 @@ int main (int argc, char **argv)
input_dictionary=(char *)Calloc(INPUT_SIZE, sizeof(char));
strncpy(input_dictionary,optarg, strlen(optarg)+1);
break;
case 'f':
if(strlen(optarg) >= INPUT_SIZE)
{
fprintf(stderr, "ERROR: Inut hash file path is bigger than %d\n", INPUT_SIZE);
exit(EXIT_FAILURE);
}
input_hash=(char *)Calloc(INPUT_SIZE, sizeof(char));
strncpy(input_hash, optarg, strlen(optarg)+1);
break;
case 'g':
gpu_id = atoi(optarg);
break;
case 'h':
usage(argv[0]);
exit(EXIT_FAILURE);
break;
case 'm':
mac_comparison = 1;
break;
case 's':
strict_check = 1;
break;
case 't':
psw_x_thread = atoi(optarg);
if(psw_x_thread <= 0)
@ -163,15 +177,6 @@ int main (int argc, char **argv)
exit(EXIT_FAILURE);
}
break;
case 'b':
gridBlocks = atoi(optarg);
break;
case 'g':
gpu_id = atoi(optarg);
break;
case 's':
strict_check = 1;
break;
default:
exit(EXIT_FAILURE);
}
@ -198,7 +203,7 @@ int main (int argc, char **argv)
}
//***********************************************************
tot_psw=(ATTACK_DEFAULT_THREADS*gridBlocks*psw_x_thread);
tot_psw=(CUDA_THREADS_NO_MAC*gridBlocks*psw_x_thread);
size_psw = tot_psw * FIXED_PASSWORD_BUFFER * sizeof(uint8_t);
//****************** GPU device *******************
if(getGPUStats())
@ -211,11 +216,21 @@ int main (int argc, char **argv)
//****************** Data from target file *******************
printf("\n====================================\nExtracting data from disk image\n====================================\n\n");
if(parse_data(input_hash, &salt, &nonce, &vmk) == BIT_FAILURE)
if(parse_data(input_hash, &salt, &nonce, &vmk, &mac) == BIT_FAILURE)
{
fprintf(stderr, "Input hash format error... exit!\n");
goto cleanup;
}
if(mac_comparison == 1 && mac == NULL)
{
fprintf(stderr, "MAC comparison option selected but no MAC string found in input hash. MAC comparison not used!\n");
mac_comparison=0;
}
else if(mac_comparison == 1)
{
tot_psw=(CUDA_THREADS_WITH_MAC*gridBlocks*psw_x_thread);
size_psw = tot_psw * FIXED_PASSWORD_BUFFER * sizeof(uint8_t);
}
//************************************************************
printf("\n\n====================================\nDictionary attack\n====================================\n\n");
@ -229,7 +244,7 @@ int main (int argc, char **argv)
//**********************************************
//************* Dictionary Attack *************
cuda_attack(input_dictionary, w_blocks_d, vmk, nonce, gridBlocks);
cuda_attack(input_dictionary, w_blocks_d, vmk, nonce, mac, gridBlocks);
//*********************************************
cleanup:

View File

@ -3,7 +3,7 @@
* 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.
* This file is part of the BitCracker project: https://github.com/e-ago/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
@ -308,12 +308,17 @@
#define ROR22(x) (((x) << 10) | ((x) >> 22))
#define ROUND(a, b, c, d, e, f, g, h, W, k) \
h += LOP3LUT_XOR(ROR6(e), ROR11(e), ROR25(e)) + LOP3LUT_XORAND(g,e,f) + UINT32_C(k) + W; \
h += LOP3LUT_XOR(ROR6(e), ROR11(e), ROR25(e)) + LOP3LUT_XORAND(g,e,f) + k + W; \
d += h; \
h += LOP3LUT_XOR(ROR2(a), ROR13(a), ROR22(a)) + LOP3LUT_ANDOR(a,b,c); /*((a & (b | c)) | (b & c)); */
h += LOP3LUT_XOR(ROR2(a), ROR13(a), ROR22(a)) + LOP3LUT_ANDOR(a,b,c);
#define ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, i, k, indexW) \
h += LOP3LUT_XOR(ROR6(e), ROR11(e), ROR25(e)) + LOP3LUT_XORAND(g,e,f) + UINT32_C(k) + tex1Dfetch(w_texture, (indexW+i)); \
h += LOP3LUT_XOR(ROR6(e), ROR11(e), ROR25(e)) + LOP3LUT_XORAND(g,e,f) + k + tex1Dfetch(w_texture, (indexW+i)); \
d += h; \
h += LOP3LUT_XOR(ROR2(a), ROR13(a), ROR22(a)) + LOP3LUT_ANDOR(a,b,c);
#define ROUND_SECOND_BLOCK_CONST(a, b, c, d, e, f, g, h, i, k, w) \
h += LOP3LUT_XOR(ROR6(e), ROR11(e), ROR25(e)) + LOP3LUT_XORAND(g,e,f) + k + w; \
d += h; \
h += LOP3LUT_XOR(ROR2(a), ROR13(a), ROR22(a)) + LOP3LUT_ANDOR(a,b,c);
@ -328,3 +333,140 @@
w_blocks_d[i] = w_blocks_d[i - 16] + w_blocks_d[i - 7] \
+ (ROR(w_blocks_d[i - 15], 7) ^ ROR(w_blocks_d[i - 15], 18) ^ (w_blocks_d[i - 15] >> 3)) \
+ (ROR(w_blocks_d[i - 2], 17) ^ ROR(w_blocks_d[i - 2], 19) ^ (w_blocks_d[i - 2] >> 10));
#define ALL_ROUND_B1_1() \
ROUND(a, b, c, d, e, f, g, h, schedule0, 0x428A2F98) \
ROUND(h, a, b, c, d, e, f, g, schedule1, 0x71374491) \
ROUND(g, h, a, b, c, d, e, f, schedule2, 0xB5C0FBCF) \
ROUND(f, g, h, a, b, c, d, e, schedule3, 0xE9B5DBA5) \
ROUND(e, f, g, h, a, b, c, d, schedule4, 0x3956C25B) \
ROUND(d, e, f, g, h, a, b, c, schedule5, 0x59F111F1) \
ROUND(c, d, e, f, g, h, a, b, schedule6, 0x923F82A4) \
ROUND(b, c, d, e, f, g, h, a, schedule7, 0xAB1C5ED5) \
ROUND(a, b, c, d, e, f, g, h, schedule8, 0xD807AA98) \
ROUND(h, a, b, c, d, e, f, g, schedule9, 0x12835B01) \
ROUND(g, h, a, b, c, d, e, f, schedule10, 0x243185BE) \
ROUND(f, g, h, a, b, c, d, e, schedule11, 0x550C7DC3) \
ROUND(e, f, g, h, a, b, c, d, schedule12, 0x72BE5D74) \
ROUND(d, e, f, g, h, a, b, c, schedule13, 0x80DEB1FE) \
ROUND(c, d, e, f, g, h, a, b, schedule14, 0x9BDC06A7) \
ROUND(b, c, d, e, f, g, h, a, schedule15, 0xC19BF174) \
ROUND(a, b, c, d, e, f, g, h, schedule16, 0xE49B69C1) \
ROUND(h, a, b, c, d, e, f, g, schedule17, 0xEFBE4786) \
ROUND(g, h, a, b, c, d, e, f, schedule18, 0x0FC19DC6) \
ROUND(f, g, h, a, b, c, d, e, schedule19, 0x240CA1CC) \
ROUND(e, f, g, h, a, b, c, d, schedule20, 0x2DE92C6F) \
ROUND(d, e, f, g, h, a, b, c, schedule21, 0x4A7484AA) \
ROUND(c, d, e, f, g, h, a, b, schedule22, 0x5CB0A9DC) \
ROUND(b, c, d, e, f, g, h, a, schedule23, 0x76F988DA) \
ROUND(a, b, c, d, e, f, g, h, schedule24, 0x983E5152) \
ROUND(h, a, b, c, d, e, f, g, schedule25, 0xA831C66D) \
ROUND(g, h, a, b, c, d, e, f, schedule26, 0xB00327C8) \
ROUND(f, g, h, a, b, c, d, e, schedule27, 0xBF597FC7) \
ROUND(e, f, g, h, a, b, c, d, schedule28, 0xC6E00BF3) \
ROUND(d, e, f, g, h, a, b, c, schedule29, 0xD5A79147) \
ROUND(c, d, e, f, g, h, a, b, schedule30, 0x06CA6351) \
ROUND(b, c, d, e, f, g, h, a, schedule31, 0x14292967)
#define ALL_ROUND_B1_2() \
ROUND(a, b, c, d, e, f, g, h, schedule0, 0x27B70A85) \
ROUND(h, a, b, c, d, e, f, g, schedule1, 0x2E1B2138) \
ROUND(g, h, a, b, c, d, e, f, schedule2, 0x4D2C6DFC) \
ROUND(f, g, h, a, b, c, d, e, schedule3, 0x53380D13) \
ROUND(e, f, g, h, a, b, c, d, schedule4, 0x650A7354) \
ROUND(d, e, f, g, h, a, b, c, schedule5, 0x766A0ABB) \
ROUND(c, d, e, f, g, h, a, b, schedule6, 0x81C2C92E) \
ROUND(b, c, d, e, f, g, h, a, schedule7, 0x92722C85) \
ROUND(a, b, c, d, e, f, g, h, schedule8, 0xA2BFE8A1) \
ROUND(h, a, b, c, d, e, f, g, schedule9, 0xA81A664B) \
ROUND(g, h, a, b, c, d, e, f, schedule10, 0xC24B8B70) \
ROUND(f, g, h, a, b, c, d, e, schedule11, 0xC76C51A3) \
ROUND(e, f, g, h, a, b, c, d, schedule12, 0xD192E819) \
ROUND(d, e, f, g, h, a, b, c, schedule13, 0xD6990624) \
ROUND(c, d, e, f, g, h, a, b, schedule14, 0xF40E3585) \
ROUND(b, c, d, e, f, g, h, a, schedule15, 0x106AA070) \
ROUND(a, b, c, d, e, f, g, h, schedule16, 0x19A4C116) \
ROUND(h, a, b, c, d, e, f, g, schedule17, 0x1E376C08) \
ROUND(g, h, a, b, c, d, e, f, schedule18, 0x2748774C) \
ROUND(f, g, h, a, b, c, d, e, schedule19, 0x34B0BCB5) \
ROUND(e, f, g, h, a, b, c, d, schedule20, 0x391C0CB3) \
ROUND(d, e, f, g, h, a, b, c, schedule21, 0x4ED8AA4A) \
ROUND(c, d, e, f, g, h, a, b, schedule22, 0x5B9CCA4F) \
ROUND(b, c, d, e, f, g, h, a, schedule23, 0x682E6FF3) \
ROUND(a, b, c, d, e, f, g, h, schedule24, 0x748F82EE) \
ROUND(h, a, b, c, d, e, f, g, schedule25, 0x78A5636F) \
ROUND(g, h, a, b, c, d, e, f, schedule26, 0x84C87814) \
ROUND(f, g, h, a, b, c, d, e, schedule27, 0x8CC70208) \
ROUND(e, f, g, h, a, b, c, d, schedule28, 0x90BEFFFA) \
ROUND(d, e, f, g, h, a, b, c, schedule29, 0xA4506CEB) \
ROUND(c, d, e, f, g, h, a, b, schedule30, 0xBEF9A3F7) \
ROUND(b, c, d, e, f, g, h, a, schedule31, 0xC67178F2)
#define ALL_ROUND_B2() \
ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 0, 0x428A2F98, 0) \
ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 1, 0x71374491, 0) \
ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 2, 0xB5C0FBCF, 0) \
ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 3, 0xE9B5DBA5, 0) \
ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 4, 0x3956C25B, indexW) \
ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 5, 0x59F111F1, indexW) \
ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 6, 0x923F82A4, indexW) \
ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 7, 0xAB1C5ED5, indexW) \
ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 8, 0xD807AA98, indexW) \
ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 9, 0x12835B01, indexW) \
ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 10, 0x243185BE, indexW) \
ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 11, 0x550C7DC3, indexW) \
ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 12, 0x72BE5D74, indexW) \
ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 13, 0x80DEB1FE, indexW) \
ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 14, 0x9BDC06A7, indexW) \
ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 15, 0xC19BF174, indexW) \
ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 16, 0xE49B69C1, indexW) \
ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 17, 0xEFBE4786, indexW) \
ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 18, 0x0FC19DC6, indexW) \
ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 19, 0x240CA1CC, indexW) \
ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 20, 0x2DE92C6F, indexW) \
ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 21, 0x4A7484AA, indexW) \
ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 22, 0x5CB0A9DC, indexW) \
ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 23, 0x76F988DA, indexW) \
ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 24, 0x983E5152, indexW) \
ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 25, 0xA831C66D, indexW) \
ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 26, 0xB00327C8, indexW) \
ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 27, 0xBF597FC7, indexW) \
ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 28, 0xC6E00BF3, indexW) \
ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 29, 0xD5A79147, indexW) \
ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 30, 0x06CA6351, indexW) \
ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 31, 0x14292967, indexW) \
ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 32, 0x27B70A85, indexW) \
ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 33, 0x2E1B2138, indexW) \
ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 34, 0x4D2C6DFC, indexW) \
ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 35, 0x53380D13, indexW) \
ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 36, 0x650A7354, indexW) \
ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 37, 0x766A0ABB, indexW) \
ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 38, 0x81C2C92E, indexW) \
ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 39, 0x92722C85, indexW) \
ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 40, 0xA2BFE8A1, indexW) \
ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 41, 0xA81A664B, indexW) \
ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 42, 0xC24B8B70, indexW) \
ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 43, 0xC76C51A3, indexW) \
ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 44, 0xD192E819, indexW) \
ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 45, 0xD6990624, indexW) \
ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 46, 0xF40E3585, indexW) \
ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 47, 0x106AA070, indexW) \
ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 48, 0x19A4C116, indexW) \
ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 49, 0x1E376C08, indexW) \
ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 50, 0x2748774C, indexW) \
ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 51, 0x34B0BCB5, indexW) \
ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 52, 0x391C0CB3, indexW) \
ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 53, 0x4ED8AA4A, indexW) \
ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 54, 0x5B9CCA4F, indexW) \
ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 55, 0x682E6FF3, indexW) \
ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 56, 0x748F82EE, indexW) \
ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 57, 0x78A5636F, indexW) \
ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 58, 0x84C87814, indexW) \
ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 59, 0x8CC70208, indexW) \
ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 60, 0x90BEFFFA, indexW) \
ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 61, 0xA4506CEB, indexW) \
ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 62, 0xBEF9A3F7, indexW) \
ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 63, 0xC67178F2, indexW)

View File

@ -3,7 +3,7 @@
* 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.
* This file is part of the BitCracker project: https://github.com/e-ago/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
@ -73,18 +73,21 @@ void print_hex(unsigned char *str, int len)
printf("%02x", str[i]);
}
int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, unsigned char ** vmk)
int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, unsigned char ** vmk, unsigned char ** mac)
{
char * hash;
char *p;
int i, salt_len, iterations, vmk_size, nonce_len;
int i, salt_size, iterations, vmk_size, nonce_size;
FILE * fphash;
char tmp[2];
int j=0;
const char zero_string[17]="0000000000000000";
(*salt) = (unsigned char *) Calloc(SALT_SIZE, sizeof(unsigned char));
(*nonce) = (unsigned char *) Calloc(NONCE_SIZE, sizeof(unsigned char));
(*vmk) = (unsigned char *) Calloc(VMK_SIZE, sizeof(unsigned char));
(*mac) = (unsigned char *) Calloc(MAC_SIZE, sizeof(unsigned char));
hash = (char *) Calloc(INPUT_HASH_SIZE, sizeof(char));
if(!input_hash)
@ -119,15 +122,15 @@ int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce,
p = strtokm(hash, "$"); // version
p = strtokm(NULL, "$"); // salt length
salt_len = atoi(p);
if(salt_len != SALT_SIZE)
salt_size = atoi(p);
if(salt_size != 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++)
for (i = 0, j = 0; i < salt_size*2; i+=2, j++)
{
tmp[0] = p[i];
tmp[1] = p[i+1];
@ -138,15 +141,15 @@ int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce,
p = strtokm(NULL, "$"); // iterations
iterations = atoi(p);
p = strtokm(NULL, "$"); // nonce length
nonce_len = atoi(p);
if(nonce_len != NONCE_SIZE)
nonce_size = atoi(p);
if(nonce_size != 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++)
for (i = 0, j = 0; i < nonce_size*2; i+=2, j++)
{
tmp[0] = p[i];
tmp[1] = p[i+1];
@ -164,20 +167,35 @@ int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce,
}
p = strtokm(NULL, "$"); // data
for (i = 0, j = 0; i < vmk_size*2; i+=2, j++)
for (i = 0, j = 0; i < MAC_SIZE*2; i+=2, j++)
{
tmp[0] = p[i];
tmp[1] = p[i+1];
long int ret = strtol(tmp, NULL, 16);
(*mac)[j] = (unsigned char)(ret); //((ARCH_INDEX(p[i * 2]) * 16) + ARCH_INDEX(p[i * 2 + 1]));
}
if(mac_comparison == 1 && !memcmp((*mac), zero_string, MAC_SIZE))
{
free(*mac);
(*mac)=NULL;
}
for (j=0; i < vmk_size*2; i+=2, j++)
{
tmp[0] = p[i];
tmp[1] = p[i+1];
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);
free(*mac);
return BIT_FAILURE;
}

View File

@ -3,7 +3,7 @@
* 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.
* This file is part of the BitCracker project: https://github.com/e-ago/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
@ -23,7 +23,7 @@
__global__ void w_block_evaluate(unsigned char salt[SALT_SIZE], int totNumIteration, unsigned char padding[40], uint32_t * w_blocks_d) { // unsigned char block[64]
uint64_t loop = (threadIdx.x+blockIdx.x*blockDim.x) + (threadIdx.y+blockIdx.y*blockDim.y) * blockDim.x * gridDim.x;
uint64_t loop = (threadIdx.x+blockIdx.x*blockDim.x);
unsigned char block[SINGLE_BLOCK_W_SIZE];
int i, j;
@ -34,10 +34,9 @@ __global__ void w_block_evaluate(unsigned char salt[SALT_SIZE], int totNumIterat
i+=8;
for(j=0; j<40; i++, j++)
{
block[i] = padding[j];
//memcpy(block, salt, SALT_SIZE);
//memcpy(block+SALT_SIZE+8, padding, 40);
}
while(loop < ITERATION_NUMBER)
{
@ -66,6 +65,7 @@ __global__ void w_block_evaluate(unsigned char salt[SALT_SIZE], int totNumIterat
LOADSCHEDULE_WPRE(13, (SINGLE_BLOCK_W_SIZE*loop)+13)
LOADSCHEDULE_WPRE(14, (SINGLE_BLOCK_W_SIZE*loop)+14)
LOADSCHEDULE_WPRE(15, (SINGLE_BLOCK_W_SIZE*loop)+15)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+16)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+17)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+18)
@ -114,8 +114,8 @@ __global__ void w_block_evaluate(unsigned char salt[SALT_SIZE], int totNumIterat
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+61)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+62)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+63)
loop = loop + (blockDim.x * gridDim.x * blockDim.y * gridDim.y);
loop += (blockDim.x * gridDim.x);
}
}

View File

@ -3,4 +3,4 @@ bitcracker_hash:
clean:
rm -rf *.o
rm -rf bitcracker_hash
rm -rf ../build/bitcracker_hash

View File

@ -70,12 +70,12 @@ void * Calloc(size_t len, size_t size) {
static int usage(char *name){
printf("\nUsage: %s -i <Encrypted memory unit> -o <output file>\n\n"
printf("\nUsage: %s -i <Image of encrypted memory unit> -o <output file>\n\n"
"Options:\n\n"
" -h, --help"
"\t\tShow this help\n"
" -i, --image"
"\t\tPath of memory unit encrypted with BitLocker\n"
"\t\tImage path of encrypted memory unit encrypted with BitLocker\n"
" -o, --outfile"
"\t\tOutput file\n\n", name);
@ -168,7 +168,38 @@ int parse_image(char * encryptedImagePath, char * outputFile)
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");
{
fprintf(stderr, "VMK encrypted with Recovery key found!\n");
#if 0
fseek(encryptedImage, 12, SEEK_CUR);
fillBuffer(encryptedImage, salt, SALT_SIZE);
fseek(encryptedImage, 83, SEEK_CUR);
char a=(unsigned char)fgetc(encryptedImage);
char b=(unsigned char)fgetc(encryptedImage);
if (( a != value_type[0]) || (b != value_type[1])) {
fprintf(stderr, "Error: VMK not encrypted with AES-CCM, a: %02x, b: %02x\n",a ,b);
match=0;
i=0;
continue;
}
else fprintf(stderr, "VMK encrypted with AES-CCM\n");
fseek(encryptedImage, 3, SEEK_CUR);
fillBuffer(encryptedImage, nonce, NONCE_SIZE);
fillBuffer(encryptedImage, mac, MAC_SIZE);
fillBuffer(encryptedImage, vmk, VMK_SIZE);
printf("Recovery key infos:\n");
printf("Salt:\n");
print_hex(salt, SALT_SIZE, stdout);
printf("\nNonce:\n");
print_hex(nonce, NONCE_SIZE, stdout);
printf("\nMac:\n");
print_hex(mac, MAC_SIZE, stdout); // hack, this should actually be entire AES-CCM encrypted block (which includes vmk)
printf("\nVMK:\n");
print_hex(vmk, VMK_SIZE, stdout);
#endif
}
else if ((c == key_protection_password[0]) && (d == key_protection_password[1]))
{
fprintf(stderr, "VMK encrypted with user password found!\n");
@ -181,11 +212,22 @@ int parse_image(char * encryptedImagePath, char * outputFile)
i=0;
continue;
}
else fprintf(stderr, "VMK encrypted with AES-CCM\n");
fseek(encryptedImage, 3, SEEK_CUR);
fillBuffer(encryptedImage, nonce, NONCE_SIZE);
fillBuffer(encryptedImage, mac, MAC_SIZE);
fillBuffer(encryptedImage, vmk, VMK_SIZE);
printf("VMK infos:\n");
printf("Salt:\n");
print_hex(salt, SALT_SIZE, stdout);
printf("\nNonce:\n");
print_hex(nonce, NONCE_SIZE, stdout);
printf("\nMac:\n");
print_hex(mac, MAC_SIZE, stdout); // hack, this should actually be entire AES-CCM encrypted block (which includes vmk)
printf("\nVMK:\n");
print_hex(vmk, VMK_SIZE, stdout);
break;
}
}
@ -203,11 +245,10 @@ int parse_image(char * encryptedImagePath, char * outputFile)
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)
printf("$%d$", VMK_SIZE + MAC_SIZE);
print_hex(mac, MAC_SIZE, stdout); // hack, this should actually be entire AES-CCM encrypted block (which includes vmk)
print_hex(vmk, VMK_SIZE, stdout);
printf("\n");
outFile = fopen(outputFile, "w");
if (!outFile) {
fprintf(stderr, "! %s : %s\n", outputFile, strerror(errno));
@ -218,8 +259,9 @@ int parse_image(char * encryptedImagePath, char * outputFile)
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)
fprintf(outFile, "$%d$", VMK_SIZE + MAC_SIZE);
print_hex(mac, MAC_SIZE, outFile);
print_hex(vmk, VMK_SIZE, outFile);
fprintf(outFile, "\n");

View File

@ -3,7 +3,7 @@
* 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.
* This file is part of the BitCracker project: https://github.com/e-ago/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

View File

@ -3,7 +3,7 @@
* 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.
* This file is part of the BitCracker project: https://github.com/e-ago/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

View File

@ -3,7 +3,7 @@
* 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.
* This file is part of the BitCracker project: https://github.com/e-ago/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

View File

@ -3,7 +3,7 @@
* 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.
* This file is part of the BitCracker project: https://github.com/e-ago/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

View File

@ -3,7 +3,7 @@
* 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.
* This file is part of the BitCracker project: https://github.com/e-ago/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

View File

@ -3,7 +3,7 @@
* 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.
* This file is part of the BitCracker project: https://github.com/e-ago/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