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

Option -m (MAC verification) in OpenCL implementation

Browse Source
This commit is contained in:
elenago 2017-11-20 20:18:47 +01:00
parent 8856b5b621
commit ef871d5603
5 changed files with 1432 additions and 141 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

17
OpenCL_version/bitcracker.h
View File

@ -58,8 +58,12 @@
#define MAC_SIZE 16
#define NONCE_SIZE 12
#define IV_SIZE 16
#define VMK_SIZE 60
#define VMK_DECRYPT_SIZE 16
#define VMK_HEADER_SIZE 12
#define VMK_BODY_SIZE 32
#define VMK_FULL_SIZE 44
#define DICT_BUFSIZE (50*1024*1024)
#define MAX_PLEN 32
@ -187,6 +191,7 @@ extern int psw_x_thread;
extern int tot_psw;
extern size_t size_psw;
extern int strict_check;
extern int mac_comparison;
extern int MAX_PASSWD_SINGLE_KERNEL;
extern int DEV_NVIDIA;
@ -207,12 +212,12 @@ extern cl_device_id* cdDevices; // OpenCL device(s)
unsigned int * w_block_precomputed(unsigned char * salt);
int readFilePassword(char ** buf, int maxNumPsw, FILE *fp);
int parse_data(char *input_hash,
unsigned char ** salt,
unsigned char ** nonce,
unsigned char ** vmk);
int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, unsigned char ** vmk, unsigned char ** mac);
char * opencl_attack(char *dname, unsigned int * w_blocks,
unsigned char * encryptedVMK,
unsigned char * nonce, unsigned char * encryptedMAC,
int gridBlocks);
char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryptedVMK, unsigned char * nonce, int gridBlocks);
void setBufferPasswordSize(size_t avail, size_t * passwordBufferSize, int * numPassword);
void * Calloc(size_t len, size_t size);

1162
OpenCL_version/kernel_attack.cl
View File

File diff suppressed because it is too large Load Diff

187
OpenCL_version/main.c
View File

@ -40,7 +40,7 @@ int psw_x_thread=8;
int tot_psw=0;
size_t size_psw=0;
size_t tot_word_mem=(SINGLE_BLOCK_SHA_SIZE * ITERATION_NUMBER * sizeof(uint32_t));
int strict_check=0;
int strict_check=0, mac_comparison=0;
// OpenCL Vars
cl_context cxGPUContext; // OpenCL context
@ -62,10 +62,12 @@ void usage(char *name)
"\t\tPath to dictionary or alphabet file\n"
" -s"
"\t\tStrict check (use only in case of false positives, faster solution)\n"
" -m"
"\t\tMAC comparison (use only in case of false positives, slower solution)\n"
" -p"
"\t\tPlatform\n"
" -g"
"\t\tGPU device number\n"
"\t\tDevice number\n"
" -t"
"\t\tSet the number of password per thread threads\n"
" -b"
@ -75,97 +77,97 @@ void usage(char *name)
int checkDeviceStatistics()
{
int i, j;
char* value;
size_t valueSize, maxWorkGroup;
cl_uint platformCount;
cl_platform_id* platforms;
cl_uint deviceCount;
cl_device_id* devices;
cl_uint maxComputeUnits, deviceAddressBits;
cl_ulong maxAllocSize, maxConstBufSize;
cl_uint ccMajor, ccMinor, registersPerBlock, warpSize, overlap;
char dname[2048];
int deviceFound=0;
char* value;
size_t valueSize, maxWorkGroup;
cl_uint platformCount;
cl_platform_id* platforms;
cl_uint deviceCount;
cl_device_id* devices;
cl_uint maxComputeUnits, deviceAddressBits;
cl_ulong maxAllocSize, maxConstBufSize;
cl_uint ccMajor, ccMinor, registersPerBlock, warpSize, overlap;
char dname[2048];
int deviceFound=0;
size_t avail, total;
// get all platforms
clGetPlatformIDs(0, NULL, &platformCount);
platforms = (cl_platform_id*) malloc(sizeof(cl_platform_id) * platformCount);
clGetPlatformIDs(platformCount, platforms, NULL);
for (i = 0; i < platformCount; i++)
{
// get all devices
clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, 0, NULL, &deviceCount);
devices = (cl_device_id*) malloc(sizeof(cl_device_id) * deviceCount);
clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, deviceCount, devices, NULL);
printf("\n# Platform: %d, # Devices: %d\n", i, deviceCount);
// for each device print critical attributes
for (j = 0; j < deviceCount; j++)
{
// get all platforms
clGetPlatformIDs(0, NULL, &platformCount);
platforms = (cl_platform_id*) malloc(sizeof(cl_platform_id) * platformCount);
clGetPlatformIDs(platformCount, platforms, NULL);
// print device name
clGetDeviceInfo(devices[j], CL_DEVICE_NAME, 0, NULL, &valueSize);
value = (char*) malloc(valueSize);
clGetDeviceInfo(devices[j], CL_DEVICE_NAME, valueSize, value, NULL);
if (platform_id == i && gpu_id == j)
{
printf("\n====================================\nSelected device: %s (ID: %d) properties\n====================================\n\n", value, j);
deviceFound=1;
}
else
printf("\n====================================\nDevice %s (ID: %d) properties\n====================================\n\n", value, j);
for (i = 0; i < platformCount; i++)
{
// get all devices
clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, 0, NULL, &deviceCount);
devices = (cl_device_id*) malloc(sizeof(cl_device_id) * deviceCount);
clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, deviceCount, devices, NULL);
free(value);
// print hardware device version
clGetDeviceInfo(devices[j], CL_DEVICE_VERSION, 0, NULL, &valueSize);
value = (char*) malloc(valueSize);
clGetDeviceInfo(devices[j], CL_DEVICE_VERSION, valueSize, value, NULL);
printf("OpenCL version supported: %s\n", value);
free(value);
// print software driver version
clGetDeviceInfo(devices[j], CL_DRIVER_VERSION, 0, NULL, &valueSize);
value = (char*) malloc(valueSize);
clGetDeviceInfo(devices[j], CL_DRIVER_VERSION, valueSize, value, NULL);
printf("Software version: %s\n", value);
free(value);
// print c version supported by compiler for device
clGetDeviceInfo(devices[j], CL_DEVICE_OPENCL_C_VERSION, 0, NULL, &valueSize);
value = (char*) malloc(valueSize);
clGetDeviceInfo(devices[j], CL_DEVICE_OPENCL_C_VERSION, valueSize, value, NULL);
printf("OpenCL C version: %s\n", value);
free(value);
clGetDeviceInfo(devices[j], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(maxAllocSize), &maxAllocSize, NULL);
printf("Max Global Memory Size: %lld\n", maxAllocSize);
GPU_MAX_GLOBAL_MEM=maxAllocSize;
maxAllocSize=0;
clGetDeviceInfo(devices[j], CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(maxAllocSize), &maxAllocSize, NULL);
printf("Max Global Memory Alloc Size: %lld\n", maxAllocSize);
GPU_MAX_MEM_ALLOC_SIZE=maxAllocSize;
printf("\n# Platform: %d, # Devices: %d\n", i, deviceCount);
// for each device print critical attributes
for (j = 0; j < deviceCount; j++)
{
clGetDeviceInfo(devices[j], CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof(maxConstBufSize), &maxConstBufSize, NULL);
printf("Max Const Memory Buffer Size: %lld\n", maxConstBufSize);
clGetDeviceInfo(devices[j], CL_DEVICE_ADDRESS_BITS, sizeof(deviceAddressBits), &deviceAddressBits, NULL);
printf("Device Address Bits: %d\n", deviceAddressBits);
// print device name
clGetDeviceInfo(devices[j], CL_DEVICE_NAME, 0, NULL, &valueSize);
value = (char*) malloc(valueSize);
clGetDeviceInfo(devices[j], CL_DEVICE_NAME, valueSize, value, NULL);
// print parallel compute units
clGetDeviceInfo(devices[j], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(maxComputeUnits), &maxComputeUnits, NULL);
printf("Parallel compute units: %d\n", maxComputeUnits);
GPU_MAX_COMPUTE_UNITS=maxComputeUnits;
if (platform_id == i && gpu_id == j)
{
printf("\n====================================\nSelected device: %s (ID: %d) properties\n====================================\n\n", value, j);
deviceFound=1;
}
else
printf("\n====================================\nDevice %s (ID: %d) properties\n====================================\n\n", value, j);
clGetDeviceInfo(devices[j], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(maxWorkGroup), &maxWorkGroup, NULL);
printf("Max Workgroup Size: %zd\n", maxWorkGroup);
GPU_MAX_WORKGROUP_SIZE=maxWorkGroup;
free(value);
// print hardware device version
clGetDeviceInfo(devices[j], CL_DEVICE_VERSION, 0, NULL, &valueSize);
value = (char*) malloc(valueSize);
clGetDeviceInfo(devices[j], CL_DEVICE_VERSION, valueSize, value, NULL);
printf("OpenCL version supported: %s\n", value);
free(value);
// print software driver version
clGetDeviceInfo(devices[j], CL_DRIVER_VERSION, 0, NULL, &valueSize);
value = (char*) malloc(valueSize);
clGetDeviceInfo(devices[j], CL_DRIVER_VERSION, valueSize, value, NULL);
printf("Software version: %s\n", value);
free(value);
// print c version supported by compiler for device
clGetDeviceInfo(devices[j], CL_DEVICE_OPENCL_C_VERSION, 0, NULL, &valueSize);
value = (char*) malloc(valueSize);
clGetDeviceInfo(devices[j], CL_DEVICE_OPENCL_C_VERSION, valueSize, value, NULL);
printf("OpenCL C version: %s\n", value);
free(value);
clGetDeviceInfo(devices[j], CL_DEVICE_VENDOR, sizeof(dname), dname, NULL);
printf("Vendor: %s\n", dname);
clGetDeviceInfo(devices[j], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(maxAllocSize), &maxAllocSize, NULL);
printf("Max Global Memory Size: %lld\n", maxAllocSize);
GPU_MAX_GLOBAL_MEM=maxAllocSize;
maxAllocSize=0;
clGetDeviceInfo(devices[j], CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(maxAllocSize), &maxAllocSize, NULL);
printf("Max Global Memory Alloc Size: %lld\n", maxAllocSize);
GPU_MAX_MEM_ALLOC_SIZE=maxAllocSize;
clGetDeviceInfo(devices[j], CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof(maxConstBufSize), &maxConstBufSize, NULL);
printf("Max Const Memory Buffer Size: %lld\n", maxConstBufSize);
clGetDeviceInfo(devices[j], CL_DEVICE_ADDRESS_BITS, sizeof(deviceAddressBits), &deviceAddressBits, NULL);
printf("Device Address Bits: %d\n", deviceAddressBits);
// print parallel compute units
clGetDeviceInfo(devices[j], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(maxComputeUnits), &maxComputeUnits, NULL);
printf("Parallel compute units: %d\n", maxComputeUnits);
GPU_MAX_COMPUTE_UNITS=maxComputeUnits;
clGetDeviceInfo(devices[j], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(maxWorkGroup), &maxWorkGroup, NULL);
printf("Max Workgroup Size: %zd\n", maxWorkGroup);
GPU_MAX_WORKGROUP_SIZE=maxWorkGroup;
clGetDeviceInfo(devices[j], CL_DEVICE_VENDOR, sizeof(dname), dname, NULL);
printf("Vendor: %s\n", dname);
if(strstr(dname, "NVIDIA") != NULL)
{
@ -297,7 +299,7 @@ int destroyClCtx()
int main (int argc, char **argv)
{
char * input_dictionary=NULL, * input_hash=NULL;
unsigned char *salt, *nonce, *vmk;
unsigned char *salt, *nonce, *vmk, *mac;
uint32_t * w_blocks_d;
long int totGlobalMem;
@ -317,7 +319,7 @@ int main (int argc, char **argv)
//*********************** Options ************************
while (1) {
opt = getopt(argc, argv, "hf:d:t:b:p:g:s");
opt = getopt(argc, argv, "hf:d:t:b:p:g:ms");
if (opt == -1)
break;
switch (opt) {
@ -357,6 +359,9 @@ int main (int argc, char **argv)
case 'g':
gpu_id = atoi(optarg);
break;
case 'm':
mac_comparison = 1;
break;
case 'p':
platform_id = atoi(optarg);
break;
@ -403,11 +408,17 @@ int main (int argc, char **argv)
//****************** Data from target file *******************
printf("\n====================================\nExtracting data from disk image\n====================================\n\n");
if(parse_data(input_hash, &salt, &nonce, &vmk) == BIT_FAILURE)
if(parse_data(input_hash, &salt, &nonce, &vmk, &mac) == BIT_FAILURE)
{
fprintf(stderr, "Input hash format error... exit!\n");
goto cleanup;
}
if(mac_comparison == 1 && mac == NULL)
{
fprintf(stderr, "MAC comparison option selected but no MAC string found in input hash. MAC comparison not used!\n");
mac_comparison=0;
}
//************************************************************
printf("\n\n====================================\nDictionary attack\n====================================\n\n");
@ -421,7 +432,7 @@ int main (int argc, char **argv)
//**********************************************
//************* Dictionary Attack *************
opencl_attack(input_dictionary, w_blocks_h, vmk, nonce, gridBlocks);
opencl_attack(input_dictionary, w_blocks_h, vmk, nonce, mac, gridBlocks);
//*********************************************
cleanup:

185
OpenCL_version/opencl_attack.c
View File

@ -45,26 +45,34 @@ static int check_match() {
}
char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryptedVMK, unsigned char * nonce, int gridBlocks)
char *opencl_attack(char *dname, unsigned int * w_blocks,
unsigned char * encryptedVMK,
unsigned char * nonce, unsigned char * encryptedMAC,
int gridBlocks)
{
cl_device_id device_id;
cl_program cpProgram;
cl_kernel ckKernelAttack;
cl_mem deviceEncryptedVMK, deviceIV, devicePassword, deviceFound, w_blocks_d;
cl_kernel ckKernelAttack;
char vmkIV[IV_SIZE], macIV[IV_SIZE], computeMacIV[IV_SIZE];
cl_mem d_vmk, d_mac, d_macIV, d_computeMacIV;
cl_mem devicePassword, deviceFound, w_blocks_d;
cl_int ciErr1, ciErr2, ccMajor;
size_t szGlobalWorkSize;
size_t szLocalWorkSize;
int numReadPassword, numPassword, passwordBufferSize, ret, totPsw=0;
char tmpIV[IV_SIZE];
FILE *fp_kernel, *fp_file_passwords;
//Very very ugly...
//Really ugly...
char fileNameAttack[] = "./OpenCL_version/kernel_attack.cl";
size_t source_size_attack, source_size;
char *source_str_attack;
size_t len = 0;
cl_int ret_cl = CL_SUCCESS, ret_cl_log = CL_SUCCESS, ret_info_kernel = CL_SUCCESS;
char optProgram[128];
unsigned int vmkIV0, vmkIV4, vmkIV8, vmkIV12;
unsigned int macIV0, macIV4, macIV8, macIV12;
unsigned int cMacIV0, cMacIV4, cMacIV8, cMacIV12;
//------- READ CL FILE ------
fp_kernel = fopen(fileNameAttack, "rb");
@ -90,19 +98,51 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp
fprintf(stderr, "crack_dict input error\n");
return NULL;
}
//-------- IV setup ------
memset(tmpIV, 0, IV_SIZE);
memcpy(tmpIV + 1, nonce, NONCE_SIZE);
if(IV_SIZE-1 - NONCE_SIZE - 1 < 0)
/*
if(tot_psw <= 0)
{
fprintf(stderr, "Nonce error\n");
fprintf(stderr, "Attack tot passwords error: %d\n", tot_psw);
return NULL;
}
*tmpIV = (unsigned char)(IV_SIZE - 1 - NONCE_SIZE - 1);
tmpIV[IV_SIZE-1] = 1;
// --------------------------------------------------------------------------
*/
//-------- vmkIV setup ------
memset(vmkIV, 0, IV_SIZE);
vmkIV[0] = (unsigned char)(IV_SIZE - 1 - NONCE_SIZE - 1);
memcpy(vmkIV + 1, nonce, NONCE_SIZE);
if(IV_SIZE-1 - NONCE_SIZE - 1 < 0)
{
fprintf(stderr, "Attack nonce error\n");
return NULL;
}
vmkIV[IV_SIZE-1] = 1;
// -----------------------
if(mac_comparison == 1)
{
//-------- macIV setup ------
memset(macIV, 0, IV_SIZE);
macIV[0] = (unsigned char)(IV_SIZE - 1 - NONCE_SIZE - 1);
memcpy(macIV + 1, nonce, NONCE_SIZE);
if(IV_SIZE-1 - NONCE_SIZE - 1 < 0)
{
fprintf(stderr, "Attack nonce error\n");
return NULL;
}
macIV[IV_SIZE-1] = 0;
// -----------------------
//-------- computeMacIV setup ------
memset(computeMacIV, 0, IV_SIZE);
computeMacIV[0] = 0x3a;
memcpy(computeMacIV + 1, nonce, NONCE_SIZE);
if(IV_SIZE-1 - NONCE_SIZE - 1 < 0)
{
fprintf(stderr, "Attack nonce error\n");
return NULL;
}
computeMacIV[IV_SIZE-1] = 0x2c;
// -----------------------
}
// ---- Open File Dictionary ----
if (!memcmp(dname, "-\0", 2)) {
fp_file_passwords= stdin;
@ -141,8 +181,16 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp
CL_ERROR(ciErr1);
}
ckKernelAttack = clCreateKernel(cpProgram, "opencl_bitcracker_attack", &ciErr1);
CL_ERROR(ciErr1);
if(mac_comparison == 1)
{
ckKernelAttack = clCreateKernel(cpProgram, "opencl_bitcracker_attack_mac", &ciErr1);
CL_ERROR(ciErr1);
}
else
{
ckKernelAttack = clCreateKernel(cpProgram, "opencl_bitcracker_attack", &ciErr1);
CL_ERROR(ciErr1);
}
size_t workgroup_size;
ret_info_kernel = clGetKernelWorkGroupInfo(ckKernelAttack, cdDevices[gpu_id], CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &workgroup_size, NULL);
@ -173,7 +221,7 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp
// --------------------------------------------------------------------------
// ------------------------------- Data setup -------------------------------
deviceEncryptedVMK = clCreateBuffer(cxGPUContext, CL_MEM_READ_WRITE, VMK_DECRYPT_SIZE*sizeof(unsigned char), NULL, &ciErr1);
d_vmk = clCreateBuffer(cxGPUContext, CL_MEM_READ_WRITE, VMK_FULL_SIZE*sizeof(unsigned char), NULL, &ciErr1);
CL_ERROR(ciErr1);
devicePassword = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, passwordBufferSize*sizeof(unsigned char), NULL, &ciErr1);
@ -184,33 +232,64 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp
w_blocks_d = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, SINGLE_BLOCK_SHA_SIZE * ITERATION_NUMBER * sizeof(unsigned int), NULL, &ciErr1);
CL_ERROR(ciErr1);
if(mac_comparison == 1)
{
d_mac = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, MAC_SIZE * sizeof(char), NULL, &ciErr1);
CL_ERROR(ciErr1);
d_macIV = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, IV_SIZE * sizeof(char), NULL, &ciErr1);
CL_ERROR(ciErr1);
d_computeMacIV = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, IV_SIZE * sizeof(char), NULL, &ciErr1);
CL_ERROR(ciErr1);
}
// --------------------------------------------------------------------------
// ------------------------------- Write buffers -------------------------------
unsigned int tmp_global = ((unsigned int *)(tmpIV))[0];
unsigned int IV0=(unsigned int )(((unsigned int )(tmp_global & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(tmp_global & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(tmp_global & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(tmp_global & 0x000000ff) << 24);
tmp_global = ((unsigned int *)(tmpIV+4))[0];
unsigned int IV4=(unsigned int )(((unsigned int )(tmp_global & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(tmp_global & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(tmp_global & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(tmp_global & 0x000000ff) << 24);
tmp_global = ((unsigned int *)(tmpIV+8))[0];
unsigned int IV8=(unsigned int )(((unsigned int )(tmp_global & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(tmp_global & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(tmp_global & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(tmp_global & 0x000000ff) << 24);
tmp_global = ((unsigned int *)(tmpIV+12))[0];
unsigned int IV12=(unsigned int )(((unsigned int )(tmp_global & 0xff000000)) >> 24) | (unsigned int )((unsigned int )(tmp_global & 0x00ff0000) >> 8) | (unsigned int )((unsigned int )(tmp_global & 0x0000ff00) << 8) | (unsigned int )((unsigned int )(tmp_global & 0x000000ff) << 24);
vmkIV0 = ((unsigned int *)(vmkIV))[0];
vmkIV4 = ((unsigned int *)(vmkIV+4))[0];
vmkIV8 = ((unsigned int *)(vmkIV+8))[0];
vmkIV12 = ((unsigned int *)(vmkIV+12))[0];
if(mac_comparison == 1)
{
macIV0 = ((unsigned int *)(macIV))[0];
macIV4 = ((unsigned int *)(macIV+4))[0];
macIV8 = ((unsigned int *)(macIV+8))[0];
macIV12 = ((unsigned int *)(macIV+12))[0];
cMacIV0 = ((unsigned int *)(computeMacIV))[0];
cMacIV4 = ((unsigned int *)(computeMacIV+4))[0];
cMacIV8 = ((unsigned int *)(computeMacIV+8))[0];
cMacIV12 = ((unsigned int *)(computeMacIV+12))[0];
}
ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, w_blocks_d, CL_TRUE, 0, SINGLE_BLOCK_SHA_SIZE * ITERATION_NUMBER * sizeof(int), w_blocks, 0, NULL, NULL);
CL_ERROR(ciErr1);
ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, deviceEncryptedVMK, CL_TRUE, 0, VMK_DECRYPT_SIZE*sizeof(char), encryptedVMK+16, 0, NULL, NULL);
ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, d_vmk, CL_TRUE, 0, VMK_FULL_SIZE*sizeof(char), encryptedVMK, 0, NULL, NULL);
CL_ERROR(ciErr1);
if(mac_comparison == 1)
{
ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, d_mac, CL_TRUE, 0, MAC_SIZE*sizeof(char), encryptedMAC, 0, NULL, NULL);
CL_ERROR(ciErr1);
ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, d_macIV, CL_TRUE, 0, IV_SIZE*sizeof(char), macIV, 0, NULL, NULL);
CL_ERROR(ciErr1);
ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, d_computeMacIV, CL_TRUE, 0, IV_SIZE*sizeof(char), computeMacIV, 0, NULL, NULL);
CL_ERROR(ciErr1);
}
// ----------------------------------------------------------------------------
szLocalWorkSize = GPU_MAX_WORKGROUP_SIZE;
szGlobalWorkSize = gridBlocks*szLocalWorkSize; //TOT THREADS
printf("Starting OpenCL attack:\n\tLocal Work Size: %d\n\tWork Group Number: %d\n\tGlobal Work Size: %zd\n\tPassword per thread: %d\n\tPassword per kernel: %d\n\tDictionary: %s\n\n",
szLocalWorkSize, gridBlocks, szGlobalWorkSize, psw_x_thread, numPassword, (fp_file_passwords == stdin)?"standard input":dname);
printf("Starting OpenCL attack:\n\tLocal Work Size: %zd\n\tWork Group Number: %d\n\tGlobal Work Size: %zd\n\tPassword per thread: %d\n\tPassword per kernel: %d\n\tDictionary: %s\n\tStrict Check (-s): %s\n\tMAC Comparison (-m): %s\n\t\n\n",
szLocalWorkSize, gridBlocks, szGlobalWorkSize, psw_x_thread, numPassword, (fp_file_passwords == stdin)?"standard input":dname, (strict_check == 1)?"Yes":"No", (mac_comparison == 1)?"Yes":"No");
int iter=0;
while(!feof(fp_file_passwords))
@ -233,23 +312,54 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp
ciErr1 |= clSetKernelArg(ckKernelAttack, 2, sizeof(cl_mem), (void*)&deviceFound);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 3, sizeof(cl_mem), (void*)&deviceEncryptedVMK);
ciErr1 |= clSetKernelArg(ckKernelAttack, 3, sizeof(cl_mem), (void*)&d_vmk);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 4, sizeof(cl_mem), (void*)&w_blocks_d);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 5, sizeof(cl_int), (void*)&IV0);
ciErr1 |= clSetKernelArg(ckKernelAttack, 5, sizeof(cl_int), (void*)&vmkIV0);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 6, sizeof(cl_int), (void*)&IV4);
ciErr1 |= clSetKernelArg(ckKernelAttack, 6, sizeof(cl_int), (void*)&vmkIV4);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 7, sizeof(cl_int), (void*)&IV8);
ciErr1 |= clSetKernelArg(ckKernelAttack, 7, sizeof(cl_int), (void*)&vmkIV8);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 8, sizeof(cl_int), (void*)&IV12);
ciErr1 |= clSetKernelArg(ckKernelAttack, 8, sizeof(cl_int), (void*)&vmkIV12);
CL_ERROR(ciErr1);
if(mac_comparison == 1)
{
ciErr1 |= clSetKernelArg(ckKernelAttack, 9, sizeof(cl_mem), (void*)&d_mac);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 10, sizeof(cl_int), (void*)&macIV0);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 11, sizeof(cl_int), (void*)&macIV4);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 12, sizeof(cl_int), (void*)&macIV8);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 13, sizeof(cl_int), (void*)&macIV12);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 14, sizeof(cl_int), (void*)&cMacIV0);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 15, sizeof(cl_int), (void*)&cMacIV4);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 16, sizeof(cl_int), (void*)&cMacIV8);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 17, sizeof(cl_int), (void*)&cMacIV12);
CL_ERROR(ciErr1);
}
// --------------------------------------------------------
time_t start,end;
@ -273,7 +383,6 @@ char *opencl_attack(char *dname, unsigned int * w_blocks, unsigned char * encryp
printf("OpenCL Kernel execution #%d\n\tEffective number psw: %d\n\tTime: %f sec\n\tPasswords x second: %10.2f pw/sec\n",
iter, numReadPassword, TIMER_ELAPSED(0)/1.0E+6, numReadPassword/(TIMER_ELAPSED(0)/1.0E+6));
// printf ("TIME timer: %d passwords in %.2lf seconds => %.2f pwd/s\n", numReadPassword, dif, (double)(numReadPassword/dif) );
ret = clFlush(cqCommandQueue);
ret = clFinish(cqCommandQueue);
@ -306,7 +415,7 @@ out:
if(cpProgram)clReleaseProgram(cpProgram);
if(w_blocks_d)clReleaseMemObject(w_blocks_d);
if(devicePassword)clReleaseMemObject(devicePassword);
if(deviceEncryptedVMK)clReleaseMemObject(deviceEncryptedVMK);
if(d_vmk)clReleaseMemObject(d_vmk);
if(deviceFound)clReleaseMemObject(deviceFound);
free(source_str_attack);

22
OpenCL_version/utils.c
View File

@ -73,7 +73,7 @@ void print_hex(unsigned char *str, int len)
printf("%02x", str[i]);
}
int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, unsigned char ** vmk)
int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, unsigned char ** vmk, unsigned char ** mac)
{
char * hash;
char *p;
@ -81,10 +81,13 @@ int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce,
FILE * fphash;
char tmp[2];
int j=0;
const char zero_string[17]="0000000000000000";
(*salt) = (unsigned char *) Calloc(SALT_SIZE, sizeof(unsigned char));
(*nonce) = (unsigned char *) Calloc(NONCE_SIZE, sizeof(unsigned char));
(*vmk) = (unsigned char *) Calloc(VMK_SIZE, sizeof(unsigned char));
(*mac) = (unsigned char *) Calloc(MAC_SIZE, sizeof(unsigned char));
hash = (char *) Calloc(INPUT_HASH_SIZE, sizeof(char));
if(!input_hash)
@ -164,7 +167,21 @@ int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce,
}
p = strtokm(NULL, "$"); // data
for (i = 0, j = 0; i < vmk_size*2; i+=2, j++)
for (i = 0, j = 0; i < MAC_SIZE*2; i+=2, j++)
{
tmp[0] = p[i];
tmp[1] = p[i+1];
long int ret = strtol(tmp, NULL, 16);
(*mac)[j] = (unsigned char)(ret); //((ARCH_INDEX(p[i * 2]) * 16) + ARCH_INDEX(p[i * 2 + 1]));
}
if(mac_comparison == 1 && !memcmp((*mac), zero_string, MAC_SIZE))
{
free(*mac);
(*mac)=NULL;
}
for (j=0; i < vmk_size*2; i+=2, j++)
{
tmp[0] = p[i];
tmp[1] = p[i+1];
@ -178,6 +195,7 @@ int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce,
free(*salt);
free(*nonce);
free(*vmk);
free(*mac);
return BIT_FAILURE;
}