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

opencl source code with Recovery Password attack

Browse Source
This commit is contained in:
elenago 2017-12-04 17:41:57 +01:00
parent 086823ed0e
commit 79e8015f23
8 changed files with 4252 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

15
src_OpenCL/Makefile Executable file
View File

@ -0,0 +1,15 @@
#!/bin/bash
FLAGS=-Wextra
NVIDIA_INCLUDE=-I/usr/local/cuda/include -L/usr/local/cuda/lib64
COMMON_INCLUDE=-I/usr/include -L/usr/lib64 -L/usr/lib
MAKE_INCLUDE=$(COMMON_INCLUDE) $(NVIDIA_INCLUDE) -cl-std=CL1.2 -O3
all:
gcc $(MAKE_INCLUDE) $(FLAGS) -o bitcracker_opencl main.c opencl_attack.c utils.c w_blocks.c -lOpenCL
# clang -framework OpenCL -o bitcracker_opencl main.c opencl_attack.c utils.c w_blocks.c
clean:
rm -rf *.o
rm -rf bitcracker_opencl

238
src_OpenCL/bitcracker.h Executable file
View File

@ -0,0 +1,238 @@
/*
* 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>
*
* 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
* 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>
#include <sys/time.h>
#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#include <CL/cl_ext.h>
#pragma OPENCL EXTENSION cl_nv_device_attribute_query : enable
#endif
#ifndef CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV
/* cl_nv_device_attribute_query extension - no extension #define since it has no functions */
#define CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV 0x4000
#define CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV 0x4001
#define CL_DEVICE_REGISTERS_PER_BLOCK_NV 0x4002
#define CL_DEVICE_WARP_SIZE_NV 0x4003
#define CL_DEVICE_GPU_OVERLAP_NV 0x4004
#define CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV 0x4005
#define CL_DEVICE_INTEGRATED_MEMORY_NV 0x4006
#endif
#define MIN(a,b) (((a)<(b))?(a):(b))
#define AUTHENTICATOR_LENGTH 16
#define AES_CTX_LENGTH 256
#define FALSE 0
#define TRUE 1
#define SALT_SIZE 16
#define MAC_SIZE 16
#define NONCE_SIZE 12
#define IV_SIZE 16
#define VMK_SIZE 60
#define VMK_HEADER_SIZE 12
#define VMK_BODY_SIZE 32
#define VMK_FULL_SIZE 44
#define RECOVERY_KEY_SIZE_CHAR 56
#define RECOVERY_PASS_BLOCKS 8
#define MODE_USER_PASS 1
#define MODE_RECV_PASS 2
#define PSW_CHAR_SIZE 64
#define PSW_INT_SIZE 16 //32 for double passwords
#define FIRST_LENGHT 27
#define SECOND_LENGHT 55
#define BLOCK_UNIT 32
#define HASH_SIZE_STRING 32
#define DICT_BUFSIZE (50*1024*1024)
#define MAX_PLEN 32
#ifndef UINT32_C
#define UINT32_C(c) c ## UL
#endif
#define HASH_SIZE 8 //32
#define ROUND_SHA_NUM 64
#define SINGLE_BLOCK_SHA_SIZE 64
#define SINGLE_BLOCK_W_SIZE 64
#define PADDING_SIZE 40
#define ITERATION_NUMBER 0x100000
#define WORD_SIZE 4
#define INPUT_SIZE 2048
#define FIXED_PART_INPUT_CHAIN_HASH 88
#define MIN_INPUT_PASSWORD_LEN 8
#define MAX_INPUT_PASSWORD_LEN 27
#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
#define BIT_FAILURE 1
#define MAX_SOURCE_SIZE (0x100000)
#define LOCAL_THREAD 768
#define MAX_NUM_PLATFORMS 10
#define MAX_DEVICE_NAME_SIZE 2048
static const char *getErrorString(cl_int error)
{
switch(error){
// run-time and JIT compiler errors
case 0: return "CL_SUCCESS";
case -1: return "CL_DEVICE_NOT_FOUND";
case -2: return "CL_DEVICE_NOT_AVAILABLE";
case -3: return "CL_COMPILER_NOT_AVAILABLE";
case -4: return "CL_MEM_OBJECT_ALLOCATION_FAILURE";
case -5: return "CL_OUT_OF_RESOURCES";
case -6: return "CL_OUT_OF_HOST_MEMORY";
case -7: return "CL_PROFILING_INFO_NOT_AVAILABLE";
case -8: return "CL_MEM_COPY_OVERLAP";
case -9: return "CL_IMAGE_FORMAT_MISMATCH";
case -10: return "CL_IMAGE_FORMAT_NOT_SUPPORTED";
case -11: return "CL_BUILD_PROGRAM_FAILURE";
case -12: return "CL_MAP_FAILURE";
case -13: return "CL_MISALIGNED_SUB_BUFFER_OFFSET";
case -14: return "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST";
case -15: return "CL_COMPILE_PROGRAM_FAILURE";
case -16: return "CL_LINKER_NOT_AVAILABLE";
case -17: return "CL_LINK_PROGRAM_FAILURE";
case -18: return "CL_DEVICE_PARTITION_FAILED";
case -19: return "CL_KERNEL_ARG_INFO_NOT_AVAILABLE";
// compile-time errors
case -30: return "CL_INVALID_VALUE";
case -31: return "CL_INVALID_DEVICE_TYPE";
case -32: return "CL_INVALID_PLATFORM";
case -33: return "CL_INVALID_DEVICE";
case -34: return "CL_INVALID_CONTEXT";
case -35: return "CL_INVALID_QUEUE_PROPERTIES";
case -36: return "CL_INVALID_COMMAND_QUEUE";
case -37: return "CL_INVALID_HOST_PTR";
case -38: return "CL_INVALID_MEM_OBJECT";
case -39: return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR";
case -40: return "CL_INVALID_IMAGE_SIZE";
case -41: return "CL_INVALID_SAMPLER";
case -42: return "CL_INVALID_BINARY";
case -43: return "CL_INVALID_BUILD_OPTIONS";
case -44: return "CL_INVALID_PROGRAM";
case -45: return "CL_INVALID_PROGRAM_EXECUTABLE";
case -46: return "CL_INVALID_KERNEL_NAME";
case -47: return "CL_INVALID_KERNEL_DEFINITION";
case -48: return "CL_INVALID_KERNEL";
case -49: return "CL_INVALID_ARG_INDEX";
case -50: return "CL_INVALID_ARG_VALUE";
case -51: return "CL_INVALID_ARG_SIZE";
case -52: return "CL_INVALID_KERNEL_ARGS";
case -53: return "CL_INVALID_WORK_DIMENSION";
case -54: return "CL_INVALID_WORK_GROUP_SIZE";
case -55: return "CL_INVALID_WORK_ITEM_SIZE";
case -56: return "CL_INVALID_GLOBAL_OFFSET";
case -57: return "CL_INVALID_EVENT_WAIT_LIST";
case -58: return "CL_INVALID_EVENT";
case -59: return "CL_INVALID_OPERATION";
case -60: return "CL_INVALID_GL_OBJECT";
case -61: return "CL_INVALID_BUFFER_SIZE";
case -62: return "CL_INVALID_MIP_LEVEL";
case -63: return "CL_INVALID_GLOBAL_WORK_SIZE";
case -64: return "CL_INVALID_PROPERTY";
case -65: return "CL_INVALID_IMAGE_DESCRIPTOR";
case -66: return "CL_INVALID_COMPILER_OPTIONS";
case -67: return "CL_INVALID_LINKER_OPTIONS";
case -68: return "CL_INVALID_DEVICE_PARTITION_COUNT";
// extension errors
case -1000: return "CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR";
case -1001: return "CL_PLATFORM_NOT_FOUND_KHR";
case -1002: return "CL_INVALID_D3D10_DEVICE_KHR";
case -1003: return "CL_INVALID_D3D10_RESOURCE_KHR";
case -1004: return "CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR";
case -1005: return "CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR";
default: return "Unknown OpenCL error";
}
}
#define CL_ERROR(errNum) \
if (errNum != CL_SUCCESS) \
{ \
fprintf(stdout, "Error in line %u in file %s: %s (%d)!!!\n\n", __LINE__, __FILE__, getErrorString(errNum), errNum); \
exit(EXIT_FAILURE); \
}
extern int gpu_id;
extern int platform_id;
extern int psw_x_thread;
extern int tot_psw;
extern size_t size_psw;
extern int strict_check;
extern int mac_comparison;
extern int attack_mode;
extern int MAX_PASSWD_SINGLE_KERNEL;
extern int DEV_NVIDIA;
extern int DEV_INTEL;
extern int DEV_AMD;
extern int CC_SM50;
//extern long int GPU_MAX_MEM_ALLOC_SIZE;
extern int GPU_MAX_COMPUTE_UNITS;
extern int GPU_MAX_WORKGROUP_SIZE;
extern long int GPU_MAX_GLOBAL_MEM;
// OpenCL Vars
extern cl_context cxGPUContext; // OpenCL context
extern cl_command_queue cqCommandQueue;// OpenCL command que
extern cl_platform_id cpPlatforms[MAX_NUM_PLATFORMS]; // OpenCL platform
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(int ** buf_i, char ** buf_c, int maxNumPsw, FILE *fp);
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);
void setBufferPasswordSize(size_t avail, size_t * passwordBufferSize, int * numPassword);
void * Calloc(size_t len, size_t size);
void print_hex(unsigned char hash[], int size);

2464
src_OpenCL/kernel_attack.cl Executable file
View File

File diff suppressed because it is too large Load Diff

162
src_OpenCL/kernel_wblocks.cl Executable file
View File

@ -0,0 +1,162 @@
/*
* 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>
*
* 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
* 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/>.
*/
#define MIN(a,b) (((a)<(b))?(a):(b))
#define AUTHENTICATOR_LENGTH 16
#define AES_CTX_LENGTH 256
#define FALSE 0
#define TRUE 1
#define SALT_SIZE 16
#define MAC_SIZE 16
#define NONCE_SIZE 12
#define IV_SIZE 16
#define VMK_SIZE 44
#define VMK_DECRYPT_SIZE 16
#define DICT_BUFSIZE (50*1024*1024)
#define MAX_PLEN 32
#define UINT32_C(c) c ## UL
#define HASH_SIZE 8 //32
#define ROUND_SHA_NUM 64
#define SINGLE_BLOCK_SHA_SIZE 64
#define SINGLE_BLOCK_W_SIZE 64
#define PADDING_SIZE 40
#define ITERATION_NUMBER 0x100000
#define WORD_SIZE 4
#define INPUT_SIZE 512
#define FIXED_PART_INPUT_CHAIN_HASH 88
#define MAX_INPUT_PASSWORD_LEN 16
#define BLOCK_UNIT 32
#define HASH_SIZE_STRING 32
#define CUDA_GRID_THREAD_X 32 /* 32 - 16*/
#define CUDA_GRID_THREAD_Y 32
#define MAX_SOURCE_SIZE (0x100000)
#define ROR(x, i) (((x) << (32 - (i))) | ((x) >> (i)))
#define LOADSCHEDULE_WPRE(i, j) \
w_blocks_d[j] = \
(unsigned int)block[i * 4 + 0] << 24 \
| (unsigned int)block[i * 4 + 1] << 16 \
| (unsigned int)block[i * 4 + 2] << 8 \
| (unsigned int)block[i * 4 + 3];
#define SCHEDULE_WPRE(i) \
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));
__kernel void opencl_bitcracker_wblocks(int totNumIteration, __global unsigned char * salt_d, __global unsigned char * padding_d, __global unsigned int * w_blocks_d)
{
unsigned long loop = get_global_id(0);
unsigned char block[SINGLE_BLOCK_W_SIZE];
int i, j;
for(i=0; i<SALT_SIZE; i++)
block[i] = salt_d[i];
i+=8;
for(j=0; j<40; i++, j++)
block[i] = padding_d[j];
while(loop < ITERATION_NUMBER)
{
block[16] = (unsigned char) (loop >> (0*8));
block[17] = (unsigned char) (loop >> (1*8));
block[18] = (unsigned char) (loop >> (2*8));
block[19] = (unsigned char) (loop >> (3*8));
block[20] = (unsigned char) (loop >> (4*8));
block[21] = (unsigned char) (loop >> (5*8));
block[22] = (unsigned char) (loop >> (6*8));
block[23] = (unsigned char) (loop >> (7*8));
LOADSCHEDULE_WPRE( 0, (SINGLE_BLOCK_W_SIZE*loop)+0)
LOADSCHEDULE_WPRE( 1, (SINGLE_BLOCK_W_SIZE*loop)+1)
LOADSCHEDULE_WPRE( 2, (SINGLE_BLOCK_W_SIZE*loop)+2)
LOADSCHEDULE_WPRE( 3, (SINGLE_BLOCK_W_SIZE*loop)+3)
LOADSCHEDULE_WPRE( 4, (SINGLE_BLOCK_W_SIZE*loop)+4)
LOADSCHEDULE_WPRE( 5, (SINGLE_BLOCK_W_SIZE*loop)+5)
LOADSCHEDULE_WPRE( 6, (SINGLE_BLOCK_W_SIZE*loop)+6)
LOADSCHEDULE_WPRE( 7, (SINGLE_BLOCK_W_SIZE*loop)+7)
LOADSCHEDULE_WPRE( 8, (SINGLE_BLOCK_W_SIZE*loop)+8)
LOADSCHEDULE_WPRE( 9, (SINGLE_BLOCK_W_SIZE*loop)+9)
LOADSCHEDULE_WPRE(10, (SINGLE_BLOCK_W_SIZE*loop)+10)
LOADSCHEDULE_WPRE(11, (SINGLE_BLOCK_W_SIZE*loop)+11)
LOADSCHEDULE_WPRE(12, (SINGLE_BLOCK_W_SIZE*loop)+12)
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)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+19)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+20)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+21)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+22)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+23)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+24)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+25)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+26)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+27)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+28)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+29)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+30)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+31)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+32)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+33)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+34)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+35)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+36)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+37)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+38)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+39)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+40)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+41)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+42)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+43)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+44)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+45)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+46)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+47)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+48)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+49)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+50)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+51)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+52)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+53)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+54)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+55)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+56)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+57)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+58)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+59)
SCHEDULE_WPRE((SINGLE_BLOCK_W_SIZE*loop)+60)
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 += get_global_size(0); //(blockDim.x * gridDim.x * blockDim.y * gridDim.y);
}
}

440
src_OpenCL/main.c Executable file
View File

@ -0,0 +1,440 @@
/*
* 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>
*
* 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
* 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 "bitcracker.h"
//#pragma OPENCL EXTENSION cl_nv_device_attribute_query : enable
int DEV_NVIDIA=0;
int DEV_INTEL=0;
int DEV_AMD=0;
int CC_SM50=0;
int MAX_PASSWD_SINGLE_KERNEL=16;
long int GPU_MAX_MEM_ALLOC_SIZE=0;
int GPU_MAX_COMPUTE_UNITS=16;
int GPU_MAX_WORKGROUP_SIZE=0;
long int GPU_MAX_GLOBAL_MEM=0;
int gpu_id=0;
int platform_id=0;
int psw_x_thread=8;
int tot_psw=0;
long int tot_word_mem=(SINGLE_BLOCK_SHA_SIZE * ITERATION_NUMBER * sizeof(uint32_t));
int strict_check=0, mac_comparison=0, attack_mode=0;
// OpenCL Vars
cl_context cxGPUContext; // OpenCL context
cl_command_queue cqCommandQueue;// OpenCL command que
cl_platform_id cpPlatforms[MAX_NUM_PLATFORMS]; // OpenCL platform
cl_uint uiNumDevices; // OpenCL total number of devices
cl_device_id* cdDevices; // OpenCL device(s)
void usage(char *name)
{
printf("\nUsage: %s -f <hash_file> -d <dictionary_file>\n\n"
"Options:\n\n"
" -h"
"\t\t\tShow this help\n"
" -f"
"\t\tPath to your input hash file (HashExtractor output)\n"
" -d"
"\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"
" -u"
"\t\tAttack User Password authentication method\n"
" -r"
"\t\tAttack Recovery Password authentication method\n"
" -p"
"\t\tPlatform\n"
" -g"
"\t\tDevice number\n"
" -t"
"\t\tSet the number of password per thread threads\n"
" -b"
"\t\tSet the number of blocks\n\n", name);
}
int checkDeviceStatistics()
{
int i, j;
char* value;
size_t valueSize, maxWorkGroup;
cl_int platformCount;
cl_platform_id* platforms;
cl_int deviceCount;
cl_device_id* devices;
cl_int maxComputeUnits, deviceAddressBits;
cl_ulong maxAllocSize, maxConstBufSize;
cl_int 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++)
{
// 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);
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;
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)
{
DEV_NVIDIA=1;
clGetDeviceInfo(devices[j], CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, sizeof(ccMajor), &ccMajor, NULL);
clGetDeviceInfo(devices[j], CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV, sizeof(ccMinor), &ccMinor, NULL);
printf("CC: %d.%d\n", ccMajor, ccMinor);
clGetDeviceInfo(devices[j], CL_DEVICE_REGISTERS_PER_BLOCK_NV, sizeof(registersPerBlock), &registersPerBlock, NULL);
printf("Registers per block: %d\n", registersPerBlock);
clGetDeviceInfo(devices[j], CL_DEVICE_WARP_SIZE_NV, sizeof(warpSize), &warpSize, NULL);
printf("Warp Size: %d\n", warpSize);
clGetDeviceInfo(devices[j], CL_DEVICE_GPU_OVERLAP_NV, sizeof(overlap), &overlap, NULL);
printf("Overlap Memory and Kernel: %d\n", overlap);
}
if (strstr(dname, "Intel") != NULL) DEV_INTEL=1;
if ((strstr(dname, "Advanced Micro") != NULL || strstr(dname, "AMD") != NULL || strstr(dname, "ATI") != NULL)) DEV_AMD=1;
if(deviceFound==1)
{
printf("\nFor this session, BitCracker requires at least %ld bytes of memory\n\n", (tot_word_mem));
if(GPU_MAX_GLOBAL_MEM < tot_word_mem)
{
fprintf(stderr, "Not enough memory available on device. Minimum required: %zd Tot memory: %ld\n", (tot_word_mem), GPU_MAX_GLOBAL_MEM);
return BIT_FAILURE;
}
if(GPU_MAX_MEM_ALLOC_SIZE < tot_word_mem)
{
fprintf(stderr, "GPU_MAX_MEM_ALLOC_SIZE: %zd Mem chunk1: %zd\n", GPU_MAX_MEM_ALLOC_SIZE, tot_word_mem);
return BIT_FAILURE;
}
break;
}
}
free(devices);
if(deviceFound == 1) break;
}
free(platforms);
if(deviceFound == 0)
{
fprintf(stderr, "Device not found! Input platform: %d, input device: %d\n", platform_id, gpu_id);
return BIT_FAILURE;
}
return BIT_SUCCESS;
}
int createClCtx()
{
cl_int clErr;
char * gpuname;
size_t gpunameSize;
// ------------------------------- OpenCL setup -------------------------------
//Get an OpenCL platform
cl_uint numPlatforms = 0;
clErr = clGetPlatformIDs(MAX_NUM_PLATFORMS, cpPlatforms, &numPlatforms);
CL_ERROR(clErr);
/* Get platform/device information */
clErr = clGetDeviceIDs(cpPlatforms[platform_id], CL_DEVICE_TYPE_ALL, 0, NULL, &uiNumDevices);
CL_ERROR(clErr);
cdDevices = (cl_device_id *)malloc(uiNumDevices * sizeof(cl_device_id) );
clErr = clGetDeviceIDs(cpPlatforms[platform_id], CL_DEVICE_TYPE_ALL, uiNumDevices, cdDevices, NULL);
CL_ERROR(clErr);
// print device name
clGetDeviceInfo(cdDevices[gpu_id], CL_DEVICE_NAME, 0, NULL, &gpunameSize);
gpuname = (char*) malloc(gpunameSize);
clGetDeviceInfo(cdDevices[gpu_id], CL_DEVICE_NAME, gpunameSize, gpuname, NULL);
printf("Setting context on Platform %d, Device '%s' (ID: %d)\n", platform_id, gpuname, gpu_id);
free(gpuname);
//Create the context
cxGPUContext = clCreateContext(0, 1, &(cdDevices[gpu_id]), NULL, NULL, &clErr);
CL_ERROR(clErr);
// Create a command-queue
cqCommandQueue = clCreateCommandQueue(cxGPUContext, cdDevices[gpu_id], 0, &clErr);
CL_ERROR(clErr);
// --------------------------------------------------------------------------
return 0;
}
int destroyClCtx()
{
if(cqCommandQueue)clReleaseCommandQueue(cqCommandQueue);
if(cxGPUContext)clReleaseContext(cxGPUContext);
return 0;
}
int main (int argc, char **argv)
{
char * input_dictionary=NULL, * input_hash=NULL;
unsigned char *salt, *nonce, *vmk, *mac;
uint32_t * w_blocks_d;
long int totGlobalMem;
//int threads = 0;
int gridBlocks = 4, ret=0, opt=0;
gpu_id=0;
platform_id=0;
printf("\n---------> BitCracker: BitLocker password cracking tool <---------\n");
if (argc < 4) {
printf("Missing argument!\n");
usage(argv[0]);
exit(EXIT_FAILURE);
}
//*********************** Options ************************
while (1) {
opt = getopt(argc, argv, "hf:d:t:b:p:g:msru");
if (opt == -1)
break;
switch (opt) {
case 'b':
gridBlocks = atoi(optarg);
break;
case 'd':
if(strlen(optarg) >= INPUT_SIZE)
{
fprintf(stderr, "ERROR: Dictionary file path is bigger than %d\n", INPUT_SIZE);
exit(EXIT_FAILURE);
}
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 'p':
platform_id = atoi(optarg);
break;
case 'r':
if(attack_mode != 0)
fprintf(stderr, "Warning: double attack type selection. Setting RECOVERY PASSWORD attack mode.\n");
attack_mode = MODE_RECV_PASS;
break;
case 's':
strict_check = 1;
break;
case 't':
psw_x_thread = atoi(optarg);
if(psw_x_thread <= 0)
{
fprintf(stderr, "ERROR: wrong password x thread number\n");
exit(EXIT_FAILURE);
}
break;
case 'u':
if(attack_mode != 0)
fprintf(stderr, "Warning: double attack type selection. Setting USER PASSWORD attack mode.\n");
attack_mode = MODE_USER_PASS;
break;
default:
exit(EXIT_FAILURE);
}
}
if (optind < argc) {
printf ("non-option ARGV-elements: ");
while (optind < argc)
printf ("%s ", argv[optind++]);
putchar ('\n');
exit(EXIT_FAILURE);
}
if (input_dictionary == NULL){
printf("Missing dictionary file!\n");
usage(argv[0]);
exit(EXIT_FAILURE);
}
if (input_hash == NULL){
printf("Missing input hash file!\n");
usage(argv[0]);
exit(EXIT_FAILURE);
}
if( attack_mode == 0 ) {
printf("\nWarning: attack type not specified (-u or -r options). Setting USER PASSWORD as default attack type\n");
attack_mode=MODE_USER_PASS;
}
//***********************************************************
if(checkDeviceStatistics())
{
fprintf(stderr, "checkDeviceStatistics error... exit!\n");
exit(EXIT_FAILURE);
}
if(createClCtx())
{
fprintf(stderr, "checkDeviceStatistics error... exit!\n");
exit(EXIT_FAILURE);
}
//****************** Data from target file *******************
printf("\n====================================\nExtracting data from disk image\n====================================\n\n");
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");
//****************** W block *******************
uint32_t * w_blocks_h = w_block_precomputed(salt);
if(!w_blocks_h)
{
fprintf(stderr, "Words error... exit!\n");
goto cleanup;
}
//**********************************************
//************* Dictionary Attack *************
opencl_attack(input_dictionary, w_blocks_h, vmk, nonce, mac, gridBlocks);
//*********************************************
cleanup:
destroyClCtx();
printf("\n\n");
return 0;
}

427
src_OpenCL/opencl_attack.c Executable file
View File

@ -0,0 +1,427 @@
/*
* 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>
*
* 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
* 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 "bitcracker.h"
#define TIMER_DEF(n) struct timeval temp_1_##n={0,0}, temp_2_##n={0,0}
#define TIMER_START(n) gettimeofday(&temp_1_##n, (struct timezone*)0)
#define TIMER_STOP(n) gettimeofday(&temp_2_##n, (struct timezone*)0)
#define TIMER_ELAPSED(n) ((temp_2_##n.tv_sec-temp_1_##n.tv_sec)*1.e6+(temp_2_##n.tv_usec-temp_1_##n.tv_usec))
int *deviceFound, *hostFound;
char *hostPassword;
int *hostPasswordInt, *devicePasswordInt;
unsigned char outPsw[PSW_CHAR_SIZE+1];
int outIndexPsw=0, match=0;
static int check_match() {
int i=0;
if (hostFound[0] >= 0){
snprintf((char*)outPsw, PSW_CHAR_SIZE+1, "%s", hostPassword+(hostFound[0]*PSW_CHAR_SIZE) );
for(i=0; i<PSW_CHAR_SIZE; i++)
if(outPsw[i] == 0x80 || outPsw[i] == 0xff) outPsw[i]='\0'; //0xffffff80
return 1;
}
return 0;
}
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;
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, tot_psw, ret, totReadPsw=0; //passwordBufferSize
FILE *fp_kernel, *fp_file_passwords;
//Really ugly...
char fileNameAttack[] = "./src_OpenCL/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");
if (!fp_kernel) {
fprintf(stderr, "Failed to load kernel.\n");
return NULL;
}
fseek(fp_kernel, 0, SEEK_END);
source_size = ftell(fp_kernel);
fseek(fp_kernel, 0, SEEK_SET);
source_str_attack = (char *)calloc(source_size+1, sizeof(char *));
source_size_attack = fread(source_str_attack, sizeof(char), source_size, fp_kernel);
if (source_size_attack != source_size)
fprintf(stderr,
"Error reading source: expected %zu, got %zu bytes.\n",
source_size, source_size_attack);
fclose(fp_kernel);
// -----------------------
if(dname == NULL || encryptedVMK == NULL || w_blocks == NULL)
{
fprintf(stderr, "crack_dict input error\n");
return NULL;
}
//-------- 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;
} else {
fp_file_passwords= fopen(dname, "r");
if (!fp_file_passwords) {
fprintf(stderr, "Cannot open file %s.\n", dname);
return NULL;
}
}
// --------------------------------------------------------------------------
// ------------------------------- Kernel setup -------------------------------
cpProgram = clCreateProgramWithSource(cxGPUContext, 1, (const char **)&source_str_attack, NULL /* (const size_t *)&source_size_attack*/, &ciErr1);
CL_ERROR(ciErr1);
clGetDeviceInfo(cdDevices[gpu_id], CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, sizeof(ccMajor), &ccMajor, NULL);
CC_SM50=0;
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 -D ATTACK_MODE=%d", CC_SM50, strict_check, attack_mode);
else
snprintf(optProgram, 128, "-I . -D DEV_NVIDIA_SM50=0 -D STRICT_CHECK=%d -D ATTACK_MODE=%d", strict_check, attack_mode);
ciErr1 = clBuildProgram(cpProgram, 1, &(cdDevices[gpu_id]), optProgram, NULL, NULL);
ret_cl = clGetProgramBuildInfo(cpProgram, cdDevices[gpu_id], CL_PROGRAM_BUILD_LOG, 0, NULL, &len);
CL_ERROR(ret_cl);
char *buffer = (char * )calloc(len+1, sizeof(char));
ret_cl_log = clGetProgramBuildInfo(cpProgram, cdDevices[gpu_id], CL_PROGRAM_BUILD_LOG, len+1, (void *)buffer, NULL);
CL_ERROR(ret_cl_log);
if(ret_cl == CL_SUCCESS && ciErr1 != CL_SUCCESS)
{
printf("Kernel Attack Build Log: \n%s\n\n", buffer);
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);
CL_ERROR(ret_info_kernel);
cl_ulong localMemSize;
ret_info_kernel = clGetKernelWorkGroupInfo(ckKernelAttack, cdDevices[gpu_id], CL_KERNEL_LOCAL_MEM_SIZE, sizeof(cl_ulong), &localMemSize, NULL);
CL_ERROR(ret_info_kernel);
size_t preferredWorkGroupSize;
ret_info_kernel = clGetKernelWorkGroupInfo(ckKernelAttack, cdDevices[gpu_id], CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, sizeof(size_t), &preferredWorkGroupSize, NULL);
CL_ERROR(ret_info_kernel);
cl_ulong privateMemSize;
ret_info_kernel = clGetKernelWorkGroupInfo(ckKernelAttack, cdDevices[gpu_id], CL_KERNEL_PRIVATE_MEM_SIZE, sizeof(cl_ulong), &privateMemSize, NULL);
CL_ERROR(ret_info_kernel);
// --------------------------------------------------------------------------
//-------- Initialize input data --------
if(GPU_MAX_WORKGROUP_SIZE > (int)workgroup_size)
GPU_MAX_WORKGROUP_SIZE = workgroup_size;
tot_psw = GPU_MAX_WORKGROUP_SIZE*gridBlocks*MAX_PASSWD_SINGLE_KERNEL;
hostPassword = (char *) Calloc(tot_psw*PSW_CHAR_SIZE, sizeof(char));
hostPasswordInt = (int *) Calloc(tot_psw*PSW_INT_SIZE, sizeof(int));
hostFound = (int *) Calloc(1, sizeof(int));
// --------------------------------------------------------------------------
// ------------------------------- Data setup -------------------------------
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, tot_psw*PSW_INT_SIZE*sizeof(unsigned int), NULL, &ciErr1);
CL_ERROR(ciErr1);
deviceFound = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, sizeof(unsigned int), NULL, &ciErr1);
CL_ERROR(ciErr1);
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 -------------------------------
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, 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;
printf("Type of attack: %s\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",
(attack_mode==MODE_USER_PASS)?"User Password":"Recovery Password", szLocalWorkSize, gridBlocks, szGlobalWorkSize, psw_x_thread, tot_psw, (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))
{
numReadPassword = readFilePassword(&hostPasswordInt, &hostPassword, tot_psw, fp_file_passwords);
ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, devicePassword, CL_TRUE, 0, tot_psw*PSW_INT_SIZE*sizeof(unsigned int), hostPasswordInt, 0, NULL, NULL);
CL_ERROR(ciErr1);
hostFound[0] = -1;
ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, deviceFound, CL_TRUE, 0, sizeof(int), hostFound, 0, NULL, NULL);
CL_ERROR(ciErr1);
ciErr1 = clSetKernelArg(ckKernelAttack, 0, sizeof(cl_int), (void*)&numReadPassword);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 1, sizeof(cl_mem), (void*)&devicePassword);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 2, sizeof(cl_mem), (void*)&deviceFound);
CL_ERROR(ciErr1);
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*)&vmkIV0);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 6, sizeof(cl_int), (void*)&vmkIV4);
CL_ERROR(ciErr1);
ciErr1 |= clSetKernelArg(ckKernelAttack, 7, sizeof(cl_int), (void*)&vmkIV8);
CL_ERROR(ciErr1);
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;
double dif;
TIMER_DEF(0);
TIMER_START(0);
time (&start);
ciErr1 = clEnqueueNDRangeKernel(cqCommandQueue, ckKernelAttack, 1, NULL, &szGlobalWorkSize, &szLocalWorkSize, 0, NULL, NULL);
CL_ERROR(ciErr1);
/* Copy result to host */
ciErr1 = clEnqueueReadBuffer(cqCommandQueue, deviceFound, CL_TRUE, 0, sizeof(unsigned int), hostFound, 0, NULL, NULL);
CL_ERROR(ciErr1);
time (&end);
TIMER_STOP(0);
dif = difftime (end,start);
printf("OpenCL Kernel execution #%d\n\tEffective number psw: %d\n\tPasswords Range:\n\t\t%s\n\t\t.....\n\t\t%s\n\tTime: %f sec\n\tPasswords x second: %10.2f pw/sec\n",
iter, numReadPassword,
(char *)(hostPassword),
(char *)(hostPassword+(PSW_CHAR_SIZE*(numReadPassword-1))),
TIMER_ELAPSED(0)/1.0E+6, numReadPassword/(TIMER_ELAPSED(0)/1.0E+6));
ret = clFlush(cqCommandQueue);
ret = clFinish(cqCommandQueue);
totReadPsw += numReadPassword;
if (hostFound[0] >= 0) {
match=check_match();
break;
}
iter++;
}
if(match==1)
printf("\n\n================================================\nOpenCL attack completed\nPasswords evaluated: %d\nPassword found: [%s]\n================================================\n\n", totReadPsw, outPsw);
else
printf("\n\n================================================\nOpenCL attack completed\nPasswords evaluated: %d\nPassword not found!\n================================================\n\n", totReadPsw);
out1:
printf("\nTot passwords evaluated: %d\n", totReadPsw);
/* Display result */
if (fp_file_passwords != stdin)
fclose(fp_file_passwords);
out:
/* Finalization */
if(ckKernelAttack)clReleaseKernel(ckKernelAttack);
if(cpProgram)clReleaseProgram(cpProgram);
if(w_blocks_d)clReleaseMemObject(w_blocks_d);
if(devicePassword)clReleaseMemObject(devicePassword);
if(d_vmk)clReleaseMemObject(d_vmk);
if(deviceFound)clReleaseMemObject(deviceFound);
free(source_str_attack);
if(match==0)
printf("Password not found\n");
return NULL;
}

328
src_OpenCL/utils.c Executable file
View File

@ -0,0 +1,328 @@
/*
* 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>
*
* 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
* 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 "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)
{
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;
}
void fillBuffer(FILE *fp, unsigned char *buffer, int size)
{
int k;
for (k = 0; k < size; k++)
buffer[k] = (unsigned char)fgetc(fp);
}
void print_hex(unsigned char *str, int len)
{
int i;
for (i = 0; i < len; ++i)
printf("%02x", str[i]);
}
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_size, iterations, vmk_size, nonce_size;
FILE * fphash;
char tmp[2];
int j=0, auth_method=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)
{
fprintf(stderr, "No input hash provided\n");
goto out;
}
fphash = fopen(input_hash, "r");
if (!fphash) {
fprintf(stderr, "! %s : %s\n", input_hash, strerror(errno));
goto out;
}
fgets(hash, INPUT_HASH_SIZE, fphash);
if(!hash)
{
fprintf(stderr, "No correct input hash provided\n");
goto out;
}
printf("Reading hash file \"%s\"\n%s", input_hash, hash);
if (strncmp(hash, HASH_TAG, HASH_TAG_LEN) != 0)
{
fprintf(stderr, "Wrong hash format\n");
goto out;
}
hash += HASH_TAG_LEN;
p = strtokm(hash, "$"); // version
auth_method = atoi(p);
if( (auth_method == 0 || auth_method == 1) && attack_mode == MODE_RECV_PASS)
{
fprintf(stderr, "Input Hash error: you choose the -r option (Recovery Password) but the input hash MUST be used with -u option (User Password).\n");
goto out;
}
else if(auth_method == 2 && attack_mode == MODE_USER_PASS)
{
fprintf(stderr, "Input Hash error: you choose the -u option (User Password) but the input hash MUST be used with -r option (Recovery Password).\n");
goto out;
}
p = strtokm(NULL, "$"); // salt length
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_size*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);
}
p = strtokm(NULL, "$"); // iterations
iterations = atoi(p);
if(iterations != ITERATION_NUMBER)
{
fprintf(stderr, "Wrong Iterations parameter\n");
goto out;
}
p = strtokm(NULL, "$"); // nonce length
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++)
{
tmp[0] = p[i];
tmp[1] = p[i+1];
long int ret = strtol(tmp, NULL, 16);
(*nonce)[j] = (unsigned char)(ret);
}
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 < 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);
}
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);
}
fclose(fphash);
return BIT_SUCCESS;
out:
fclose(fphash);
free(*salt);
free(*nonce);
free(*vmk);
free(*mac);
return BIT_FAILURE;
}
int readFilePassword(int ** buf_i, char ** buf_c, int maxNumPsw, FILE *fp) {
int i=0, j=0, k=0, size=0, count=0;
char tmp[PSW_CHAR_SIZE], tmp2[PSW_CHAR_SIZE], *p;
memset(tmp, 0, PSW_CHAR_SIZE);
fseek(fp, 0, SEEK_SET);
if (fp == NULL || feof(fp) || buf_i == NULL)
return -1;
while(fgets(tmp, PSW_CHAR_SIZE, fp) && (i < maxNumPsw)) {
size = (strlen(tmp)-1);
j=0; k=0; count=0;
if(tmp[0] == '\n' || size < MIN_INPUT_PASSWORD_LEN || size > SECOND_LENGHT) continue;
memcpy(( (*buf_c)+(i*PSW_CHAR_SIZE)), tmp, size);
//Recovery password
if(attack_mode == MODE_RECV_PASS)
{
memset(tmp2, 0, PSW_CHAR_SIZE);
p = strtokm(tmp, "-");
do
{
//Dislocker, Recovery Password checks
if( ((atoi(p) % 11) != 0) || (atoi(p) >= 720896) ) break;
int8_t check_digit = (int8_t) ( p[0] - p[1] + p[2] - p[3] + p[4] - 48 ) % 11;
if( check_digit < 0 ) check_digit = (int8_t) check_digit + 11;
if( check_digit != (p[5] - 48)) break;
((uint16_t*)(tmp2+count))[0] = (uint16_t)(atoi(p) / 11);
p = strtokm(NULL, "-");
count+=2;
} while(p != NULL);
if(count != (RECOVERY_PASS_BLOCKS*2)) continue;
((*buf_i)+(i*PSW_INT_SIZE))[0] = ( (((unsigned int)tmp2[0] ) << 24) & 0xFF000000) |
( (((unsigned int)tmp2[0+1]) << 16) & 0x00FF0000) |
( (((unsigned int)tmp2[0+2]) << 8) & 0x0000FF00) |
( (((unsigned int)tmp2[0+3]) << 0) & 0x000000FF);
((*buf_i)+(i*PSW_INT_SIZE))[1] = ( (((unsigned int)tmp2[4]) << 24) & 0xFF000000) |
( (((unsigned int)tmp2[4+1]) << 16) & 0x00FF0000) |
( (((unsigned int)tmp2[4+2]) << 8) & 0x0000FF00) |
( (((unsigned int)tmp2[4+3]) << 0) & 0x000000FF);
((*buf_i)+(i*PSW_INT_SIZE))[2] = ( (((unsigned int)tmp2[8]) << 24) & 0xFF000000) |
( (((unsigned int)tmp2[8+1]) << 16) & 0x00FF0000) |
( (((unsigned int)tmp2[8+2]) << 8) & 0x0000FF00) |
( (((unsigned int)tmp2[8+3]) << 0) & 0x000000FF);
((*buf_i)+(i*PSW_INT_SIZE))[3] = ( (((unsigned int)tmp2[12]) << 24) & 0xFF000000) |
( (((unsigned int)tmp2[12+1]) << 16) & 0x00FF0000) |
( (((unsigned int)tmp2[12+2]) << 8) & 0x0000FF00) |
( (((unsigned int)tmp2[12+3]) << 0) & 0x000000FF);
((*buf_i)+(i*PSW_INT_SIZE))[4] = 0x80000000;
((*buf_i)+(i*PSW_INT_SIZE))[5] = 0;
((*buf_i)+(i*PSW_INT_SIZE))[6] = 0;
((*buf_i)+(i*PSW_INT_SIZE))[7] = 0;
((*buf_i)+(i*PSW_INT_SIZE))[8] = 0;
((*buf_i)+(i*PSW_INT_SIZE))[9] = 0;
((*buf_i)+(i*PSW_INT_SIZE))[10] = 0;
((*buf_i)+(i*PSW_INT_SIZE))[11] = 0;
((*buf_i)+(i*PSW_INT_SIZE))[12] = 0;
((*buf_i)+(i*PSW_INT_SIZE))[13] = 0;
((*buf_i)+(i*PSW_INT_SIZE))[14] = 0;
((*buf_i)+(i*PSW_INT_SIZE))[15] = 0x80;
}
else //User Password
{
tmp[size] = 0x80;
do
{
((*buf_i)+(i*PSW_INT_SIZE)+j)[0] = ( (((unsigned int)tmp[k]) << 24) & 0xFF000000);
k++;
if(k <= size)
((*buf_i)+(i*PSW_INT_SIZE)+j)[0] = ((*buf_i)+(i*PSW_INT_SIZE)+j)[0] | ( (((unsigned int)tmp[k]) << 8) & 0x0000FF00);
j++;
k++;
} while(k <= size);
if(size <= FIRST_LENGHT)
{
((*buf_i)+(i*PSW_INT_SIZE)+14)[0] = 0;
((*buf_i)+(i*PSW_INT_SIZE)+15)[0] = ((int)(((size*2) << 3) >> 8)) << 8 | ((int)((size*2) << 3));
}
else
{
// Next release!
fprintf(stderr, "ERROR!\n");
exit(EXIT_FAILURE);
((*buf_i)+(i*PSW_INT_SIZE)+30)[0] = 0;
((*buf_i)+(i*PSW_INT_SIZE)+31)[0] = ((uint8_t)(((size*2) << 3) >> 8)) << 8 | ((uint8_t)((size*2) << 3));
}
}
memset(tmp, 0, PSW_CHAR_SIZE);
i++;
}
return i;
}

178
src_OpenCL/w_blocks.c Executable file
View File

@ -0,0 +1,178 @@
/*
* 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>
*
* 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 "bitcracker.h"
unsigned int * w_block_precomputed(unsigned char * salt)
{
cl_device_id device_id;
cl_program cpProgram;
cl_kernel ckKernelWBlocks;
cl_mem salt_d, padding_d, w_blocks_d;
cl_int ciErr1;
size_t szGlobalWorkSize;
size_t szLocalWorkSize;
int i;
FILE *fp_kernel;
time_t start,end;
double dif;
//Very very ugly...
const char fileNameWBlocks[] = "./src_OpenCL/kernel_wblocks.cl";
size_t source_size_wbocks;
char *source_str_wbocks;
unsigned char * padding;
uint64_t msgLen;
unsigned int * w_blocks_h;
int iter_num;
size_t len = 0;
cl_int ret = CL_SUCCESS, ret_cl_log = CL_SUCCESS, ret_info_kernel = CL_SUCCESS;
if(salt == NULL)
return FALSE;
//------- READ CL FILE ------
fp_kernel = fopen(fileNameWBlocks, "rb");
if (!fp_kernel) {
fprintf(stderr, "Failed to load kernel.\n");
return NULL;
}
source_str_wbocks = (char *)malloc(MAX_SOURCE_SIZE);
source_size_wbocks = fread(source_str_wbocks, 1, MAX_SOURCE_SIZE, fp_kernel);
fclose(fp_kernel);
// -----------------------
padding = (unsigned char *) Calloc(PADDING_SIZE, sizeof(unsigned char));
padding[0] = 0x80;
memset(padding+1, 0, 31);
msgLen = (FIXED_PART_INPUT_CHAIN_HASH << 3);
for (i = 0; i < 8; i++)
padding[PADDING_SIZE-1-i] = (uint8_t)(msgLen >> (i * 8));
// ------------------------------- Data setup -------------------------------
salt_d = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, SALT_SIZE * sizeof(char), NULL, &ciErr1);
CL_ERROR(ciErr1);
padding_d = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, PADDING_SIZE * sizeof(char), NULL, &ciErr1);
CL_ERROR(ciErr1);
w_blocks_d = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, SINGLE_BLOCK_SHA_SIZE * ITERATION_NUMBER * sizeof(int), NULL, &ciErr1);
CL_ERROR(ciErr1);
w_blocks_h = (unsigned int *) Calloc((SINGLE_BLOCK_SHA_SIZE*ITERATION_NUMBER), sizeof(int));
if(!w_blocks_h)
goto out;
// --------------------------------------------------------------------------
// ------------------------------- Kernel setup -------------------------------
cpProgram = clCreateProgramWithSource(cxGPUContext, 1, (const char **)&source_str_wbocks, (const size_t *)&source_size_wbocks, &ciErr1);
CL_ERROR(ciErr1);
ciErr1 = clBuildProgram(cpProgram, 1, &(cdDevices[gpu_id]), "-I .", NULL, NULL);
ret = clGetProgramBuildInfo(cpProgram, cdDevices[gpu_id], CL_PROGRAM_BUILD_LOG, 0, NULL, &len);
CL_ERROR(ret);
char *buffer = calloc(len, sizeof(char));
ret_cl_log = clGetProgramBuildInfo(cpProgram, cdDevices[gpu_id], CL_PROGRAM_BUILD_LOG, len, buffer, NULL);
CL_ERROR(ret_cl_log);
if(ret_cl_log == CL_SUCCESS && ciErr1 != CL_SUCCESS)
{
printf("Kernel Blocks Build Log: \n%s\n\n", buffer);
CL_ERROR(ciErr1);
}
ckKernelWBlocks = clCreateKernel(cpProgram, "opencl_bitcracker_wblocks", &ciErr1);
CL_ERROR(ciErr1);
// --------------------------------------------------------------------------
// ------------------------------- Write static buffers -------------------------------
ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, salt_d, CL_TRUE, 0, SALT_SIZE * sizeof(unsigned char), salt, 0, NULL, NULL);
CL_ERROR(ciErr1);
ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, padding_d, CL_TRUE, 0, PADDING_SIZE * sizeof(unsigned char), padding, 0, NULL, NULL);
CL_ERROR(ciErr1);
// --------------------------------------------------------------------------
iter_num = ITERATION_NUMBER;
ciErr1 = clSetKernelArg(ckKernelWBlocks, 0, sizeof(cl_int), (void*)&iter_num);
CL_ERROR(ciErr1);
ciErr1 = clSetKernelArg(ckKernelWBlocks, 1, sizeof(cl_mem), (void*)&salt_d);
CL_ERROR(ciErr1);
ciErr1 = clSetKernelArg(ckKernelWBlocks, 2, sizeof(cl_mem), (void*)&padding_d);
CL_ERROR(ciErr1);
ciErr1 = clSetKernelArg(ckKernelWBlocks, 3, sizeof(cl_mem), (void*)&w_blocks_d);
CL_ERROR(ciErr1);
// ---------------------- Launch kernel
size_t workgroup_size;
ret_info_kernel = clGetKernelWorkGroupInfo(ckKernelWBlocks, cdDevices[gpu_id], CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &workgroup_size, NULL);
CL_ERROR(ret_info_kernel);
cl_ulong localMemSize;
ret_info_kernel = clGetKernelWorkGroupInfo(ckKernelWBlocks, cdDevices[gpu_id], CL_KERNEL_LOCAL_MEM_SIZE, sizeof(cl_ulong), &localMemSize, NULL);
CL_ERROR(ret_info_kernel);
size_t preferredWorkGroupSize;
ret_info_kernel = clGetKernelWorkGroupInfo(ckKernelWBlocks, cdDevices[gpu_id], CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, sizeof(size_t), &preferredWorkGroupSize, NULL);
CL_ERROR(ret_info_kernel);
cl_ulong privateMemSize;
ret_info_kernel = clGetKernelWorkGroupInfo(ckKernelWBlocks, cdDevices[gpu_id], CL_KERNEL_PRIVATE_MEM_SIZE, sizeof(cl_ulong), &privateMemSize, NULL);
CL_ERROR(ret_info_kernel);
// --------------------------------------------------------------------------
//-------- Initialize input data --------
szLocalWorkSize = workgroup_size;
szGlobalWorkSize = 16*szLocalWorkSize;
time (&start);
ciErr1 = clEnqueueNDRangeKernel(cqCommandQueue, ckKernelWBlocks, 1, NULL, &szGlobalWorkSize, &szLocalWorkSize, 0, NULL, NULL);
CL_ERROR(ciErr1);
ciErr1 = clEnqueueReadBuffer(cqCommandQueue, w_blocks_d, CL_TRUE, 0, SINGLE_BLOCK_SHA_SIZE*ITERATION_NUMBER*sizeof(unsigned int), w_blocks_h, 0, NULL, NULL);
CL_ERROR(ciErr1);
time (&end);
dif = difftime (end,start);
ret = clFlush(cqCommandQueue);
ret = clFinish(cqCommandQueue);
out:
if(ckKernelWBlocks)clReleaseKernel(ckKernelWBlocks);
if(cpProgram)clReleaseProgram(cpProgram);
clReleaseMemObject(salt_d);
clReleaseMemObject(padding_d);
clReleaseMemObject(w_blocks_d);
free(padding);
free(source_str_wbocks);
/*
* Useless
* unsigned int * w_blocks_h = NULL;
* w_blocks_h = (unsigned int *) Calloc((SINGLE_BLOCK_SHA_SIZE*ITERATION_NUMBER), sizeof(unsigned int));
* BITCRACKER_CUDA_CHECK( cudaMemcpy(w_blocks_h, w_blocks_d, SINGLE_BLOCK_SHA_SIZE * ITERATION_NUMBER * sizeof(unsigned int), cudaMemcpyDeviceToHost) );
*/
return w_blocks_h;
}