diff --git a/CUDA_version/Makefile b/CUDA_version/Makefile index 12da431..c5c3f13 100755 --- a/CUDA_version/Makefile +++ b/CUDA_version/Makefile @@ -3,4 +3,4 @@ bitcracker_cuda: clean: rm -rf *.o - rm -rf bitcracker_cuda + rm -rf ../build/bitcracker_cuda diff --git a/CUDA_version/aes_header.h b/CUDA_version/aes_header.h index ed2be23..560d739 100755 --- a/CUDA_version/aes_header.h +++ b/CUDA_version/aes_header.h @@ -3,7 +3,7 @@ * Copyright (C) 2013-2017 Elena Ago * Massimo Bernaschi * - * 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 diff --git a/CUDA_version/bitcracker.h b/CUDA_version/bitcracker.h index 54d15fa..7ff953a 100755 --- a/CUDA_version/bitcracker.h +++ b/CUDA_version/bitcracker.h @@ -3,7 +3,7 @@ * Copyright (C) 2013-2017 Elena Ago * Massimo Bernaschi * - * 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 #include #include +#include #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); diff --git a/CUDA_version/cuda_attack.cu b/CUDA_version/cuda_attack.cu index cdad1f2..f7797f6 100755 --- a/CUDA_version/cuda_attack.cu +++ b/CUDA_version/cuda_attack.cu @@ -3,7 +3,7 @@ * Copyright (C) 2013-2017 Elena Ago * Massimo Bernaschi * - * 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 @@ -43,18 +43,23 @@ static int check_match(int iStream) { return 0; } -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) +{ FILE *fp; - int indexStream, numReadPassword[2], firstLoop, match=0; + int indexStream, numReadPassword[2], firstLoop, match=0; long long totPsw = 0; - - uint8_t tmpIV[IV_SIZE], *deviceEncryptedVMK, *deviceIV; - - + uint8_t vmkIV[IV_SIZE], *d_vmkIV, *d_vmk; + uint8_t macIV[IV_SIZE], *d_macIV, *d_mac; + uint8_t computeMacIV[IV_SIZE], *d_computeMacIV; + int cudaThreads=CUDA_THREADS_NO_MAC; cudaEvent_t start[2], stop[2]; - cudaStream_t stream[2]; - float elapsedTime; + cudaStream_t stream[2]; + float elapsedTime; + int w_blocks_h[4]; if(dname == NULL || w_blocks_d == NULL || encryptedVMK == NULL) { @@ -68,18 +73,47 @@ char *cuda_attack(char *dname, uint32_t * w_blocks_d, unsigned char * encryptedV return NULL; } - //-------- IV setup ------ - memset(tmpIV, 0, IV_SIZE); - memcpy(tmpIV + 1, nonce, NONCE_SIZE); + //-------- 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; } - *tmpIV = (unsigned char)(IV_SIZE - 1 - NONCE_SIZE - 1); - tmpIV[IV_SIZE-1] = 1; + vmkIV[IV_SIZE-1] = 1; // ----------------------- + if(mac_comparison == 1) + { + cudaThreads=CUDA_THREADS_WITH_MAC; + + //-------- 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 = stdin; @@ -101,11 +135,37 @@ char *cuda_attack(char *dname, uint32_t * w_blocks_d, unsigned char * encryptedV // ------------------------ // ---- CUDA VARIABLES ---- - BITCRACKER_CUDA_CHECK( cudaMalloc( (void ** ) &deviceEncryptedVMK, VMK_DECRYPT_SIZE*sizeof(uint8_t)) ); - BITCRACKER_CUDA_CHECK( cudaMemcpy(deviceEncryptedVMK, (encryptedVMK+16), VMK_DECRYPT_SIZE*sizeof(uint8_t), cudaMemcpyHostToDevice) ); - BITCRACKER_CUDA_CHECK( cudaMalloc( (void ** ) &deviceIV, IV_SIZE*sizeof(uint8_t)) ); - BITCRACKER_CUDA_CHECK( cudaMemcpy(deviceIV, tmpIV, IV_SIZE*sizeof(uint8_t), cudaMemcpyHostToDevice) ); + BITCRACKER_CUDA_CHECK( cudaMalloc( (void ** ) &d_vmk, VMK_FULL_SIZE*sizeof(uint8_t)) ); + BITCRACKER_CUDA_CHECK( cudaMemcpy(d_vmk, (encryptedVMK), VMK_FULL_SIZE*sizeof(uint8_t), cudaMemcpyHostToDevice) ); + + BITCRACKER_CUDA_CHECK( cudaMalloc( (void ** ) &d_vmkIV, IV_SIZE*sizeof(uint8_t)) ); + BITCRACKER_CUDA_CHECK( cudaMemcpy(d_vmkIV, vmkIV, IV_SIZE*sizeof(uint8_t), cudaMemcpyHostToDevice) ); + printf("\nVMK IV: "); + print_hex(vmkIV, IV_SIZE); + printf("\n"); + + if(mac_comparison == 1) + { + BITCRACKER_CUDA_CHECK( cudaMalloc( (void ** ) &d_mac, MAC_SIZE*sizeof(uint8_t)) ); + BITCRACKER_CUDA_CHECK( cudaMemcpy(d_mac, encryptedMAC, MAC_SIZE*sizeof(uint8_t), cudaMemcpyHostToDevice) ); + + BITCRACKER_CUDA_CHECK( cudaMalloc( (void ** ) &d_macIV, IV_SIZE*sizeof(uint8_t)) ); + BITCRACKER_CUDA_CHECK( cudaMemcpy(d_macIV, macIV, IV_SIZE*sizeof(uint8_t), cudaMemcpyHostToDevice) ); + + BITCRACKER_CUDA_CHECK( cudaMalloc( (void ** ) &d_computeMacIV, IV_SIZE*sizeof(uint8_t)) ); + BITCRACKER_CUDA_CHECK( cudaMemcpy(d_computeMacIV, computeMacIV, IV_SIZE*sizeof(uint8_t), cudaMemcpyHostToDevice) ); + + printf("VMK Body: "); + print_hex((encryptedVMK), VMK_FULL_SIZE); + printf("\nMAC Encrypted: "); + print_hex(encryptedMAC, MAC_SIZE); + printf("\nMAC IV: "); + print_hex(macIV, IV_SIZE); + printf("\nCompute MAC IV: "); + print_hex(computeMacIV, IV_SIZE); + printf("\n\n"); + } BITCRACKER_CUDA_CHECK( cudaMalloc( (void ** ) &devicePassword[0], (size_psw * sizeof(uint8_t)) ) ); BITCRACKER_CUDA_CHECK( cudaMalloc( (void ** ) &devicePassword[1], (size_psw * sizeof(uint8_t)) ) ); @@ -125,16 +185,23 @@ char *cuda_attack(char *dname, uint32_t * w_blocks_d, unsigned char * encryptedV BITCRACKER_CUDA_CHECK( cudaEventCreate(&stop[1]) ); // --------------------- + BITCRACKER_CUDA_CHECK( cudaMemcpy(w_blocks_h, w_blocks_d, 4*sizeof(int), cudaMemcpyDeviceToHost) ); + // -------- TEXTURE -------- BITCRACKER_CUDA_CHECK(cudaBindTexture(NULL, w_texture, w_blocks_d, (SINGLE_BLOCK_SHA_SIZE * ITERATION_NUMBER * sizeof(uint32_t)))); BITCRACKER_CUDA_CHECK(cudaBindTexture(NULL, w_password0, devicePassword[0], (size_psw * sizeof(uint8_t)))); BITCRACKER_CUDA_CHECK(cudaBindTexture(NULL, w_password1, devicePassword[1], (size_psw * sizeof(uint8_t)))); // ------------------------- - // BITCRACKER_CUDA_CHECK (cudaDeviceSetCacheConfig( cudaFuncCachePreferL1 ) ); + BITCRACKER_CUDA_CHECK (cudaDeviceSetCacheConfig( cudaFuncCachePreferL1 ) ); - printf("Starting CUDA attack:\n\tCUDA Threads: %d\n\tCUDA Blocks: %d\n\tPsw per thread: %d\n\tMax Psw per kernel: %d\n\tDictionary: %s\n\n", - ATTACK_DEFAULT_THREADS, gridBlocks, psw_x_thread, tot_psw, (fp == stdin)?"standard input":dname); + printf("Starting CUDA attack:\n\tCUDA Threads: %d\n\tCUDA Blocks: %d\n\tPsw per thread: %d\n\tMax Psw per kernel: %d\n\tDictionary: %s\n\tStrict Check (-s): %d\n\tMAC Comparison (-m): %d\n\t\n\n", + cudaThreads, gridBlocks, psw_x_thread, tot_psw, (fp == stdin)?"standard input":dname, (strict_check == 1)?1:0, (mac_comparison == 1)?1:0); + + uint32_t s0 = ((uint32_t)salt[0] ) << 24 | ((uint32_t)salt[1] ) << 16 | ((uint32_t)salt[2] ) << 8 | ((uint32_t)salt[3]); + uint32_t s1 = ((uint32_t)salt[4] ) << 24 | ((uint32_t)salt[5] ) << 16 | ((uint32_t)salt[6] ) << 8 | ((uint32_t)salt[7]); + uint32_t s2 = ((uint32_t)salt[8] ) << 24 | ((uint32_t)salt[9] ) << 16 | ((uint32_t)salt[10]) << 8 | ((uint32_t)salt[11]); + uint32_t s3 = ((uint32_t)salt[12]) << 24 | ((uint32_t)salt[13]) << 16 | ((uint32_t)salt[14]) << 8 | ((uint32_t)salt[15]); indexStream = 1; firstLoop=TRUE; @@ -144,11 +211,28 @@ char *cuda_attack(char *dname, uint32_t * w_blocks_d, unsigned char * encryptedV BITCRACKER_CUDA_CHECK( cudaMemcpyAsync(devicePassword[indexStream], hostPassword[indexStream], size_psw, cudaMemcpyHostToDevice, stream[indexStream]) ); - if(firstLoop == FALSE) - BITCRACKER_CUDA_CHECK( cudaStreamSynchronize(stream[indexStream^1]) ); + if(firstLoop == FALSE) BITCRACKER_CUDA_CHECK( cudaStreamSynchronize(stream[indexStream^1]) ); BITCRACKER_CUDA_CHECK( cudaEventRecord(start[indexStream], stream[indexStream]) ); - decrypt_vmk<<>>(indexStream, numReadPassword[indexStream], deviceFound[indexStream], deviceEncryptedVMK, deviceIV, strict_check); + if(mac_comparison == 1) + { + printf("launch attack with mac\n"); + decrypt_vmk_with_mac<<>>(indexStream, + numReadPassword[indexStream], deviceFound[indexStream], d_vmk, d_vmkIV, d_mac, d_macIV, d_computeMacIV, + w_blocks_h[0], w_blocks_h[1], w_blocks_h[2], w_blocks_h[3], + s0, s1, s2, s3); + } + else + { + printf("launch attack normal\n"); + decrypt_vmk<<>>(indexStream, + numReadPassword[indexStream], deviceFound[indexStream], d_vmk, d_vmkIV, strict_check, + w_blocks_h[0], w_blocks_h[1], w_blocks_h[2], w_blocks_h[3], + s0, s1, s2, s3); + } + + + BITCRACKER_CUDA_CHECK_LAST_ERROR(); BITCRACKER_CUDA_CHECK( cudaEventRecord(stop[indexStream], stream[indexStream]) ); BITCRACKER_CUDA_CHECK( cudaMemcpyAsync(hostFound[indexStream], deviceFound[indexStream], sizeof(unsigned int), cudaMemcpyDeviceToHost, stream[indexStream]) ); @@ -159,13 +243,13 @@ char *cuda_attack(char *dname, uint32_t * w_blocks_d, unsigned char * encryptedV BITCRACKER_CUDA_CHECK( cudaEventElapsedTime(&elapsedTime, start[indexStream^1], stop[indexStream^1]) ); printf("CUDA Kernel execution:\n\tStream %d\n\tEffective number psw: %d\n\tTime: %f sec\n\tPasswords x second: %8.2f pw/sec\n", - indexStream^1, numReadPassword[indexStream^1], ATTACK_DEFAULT_THREADS, gridBlocks, (elapsedTime/1000.0), numReadPassword[indexStream^1]/(elapsedTime/1000.0)); + indexStream^1, numReadPassword[indexStream^1], cudaThreads, gridBlocks, (elapsedTime/1000.0), numReadPassword[indexStream^1]/(elapsedTime/1000.0)); match=check_match(indexStream^1); if(match) break; } - firstLoop = FALSE; + firstLoop = FALSE; } BITCRACKER_CUDA_CHECK( cudaStreamSynchronize(stream[indexStream]) ); @@ -206,7 +290,11 @@ char *cuda_attack(char *dname, uint32_t * w_blocks_d, unsigned char * encryptedV #define END_STRING 0x80 //0xFF -__global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsigned char * vmkKey, unsigned char * IV, int strict_check) { +__global__ 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 + ) +{ uint32_t schedule0, schedule1, schedule2, schedule3, schedule4, schedule5, schedule6, schedule7, schedule8, schedule9; uint32_t schedule10, schedule11, schedule12, schedule13, schedule14, schedule15, schedule16, schedule17, schedule18, schedule19; @@ -216,7 +304,7 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig uint32_t hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7; uint32_t a, b, c, d, e, f, g, h; - int gIndex = (threadIdx.x+blockIdx.x*blockDim.x); + int gIndex = (threadIdx.x+blockIdx.x*blockDim.x); int index_generic; int indexW=(gIndex*FIXED_PASSWORD_BUFFER); int8_t curr_fetch=0; @@ -375,75 +463,10 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig schedule15 = ((uint8_t)((index_generic << 3) >> 8)) << 8 | ((uint8_t)(index_generic << 3)); ALL_SCHEDULE_LAST16() - - 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) - + ALL_ROUND_B1_1() ALL_SCHEDULE32() - - 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) - + ALL_ROUND_B1_2() + first_hash0 += a; first_hash1 += b; first_hash2 += c; @@ -453,8 +476,6 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig first_hash6 += g; first_hash7 += h; - - //----------------------------------------------------- SECOND HASH ------------------------------------------------ //old loadschedule schedule0 = first_hash0; @@ -493,74 +514,9 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig h = first_hash7; ALL_SCHEDULE_LAST16() - - 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) - + ALL_ROUND_B1_1() ALL_SCHEDULE32() - - 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) + ALL_ROUND_B1_2() first_hash0 += a; first_hash1 += b; @@ -571,7 +527,7 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig first_hash6 += g; first_hash7 += h; - //----------------------------------------------------- LOOP HASH ------------------------------------------------ +//----------------------------------------------------- LOOP HASH ------------------------------------------------ hash0=0; hash1=0; @@ -614,74 +570,9 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig schedule15 = first_hash7; ALL_SCHEDULE_LAST16() - - 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) - + ALL_ROUND_B1_1() ALL_SCHEDULE32() - - 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) + ALL_ROUND_B1_2() hash0 = UINT32_C(0x6A09E667) + a; hash1 = UINT32_C(0xBB67AE85) + b; @@ -701,10 +592,10 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig g = hash6; h = hash7; - 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_CONST(a, b, c, d, e, f, g, h, 0, 0x428A2F98, v0) + ROUND_SECOND_BLOCK_CONST(h, a, b, c, d, e, f, g, 1, 0x71374491, v1) + ROUND_SECOND_BLOCK_CONST(g, h, a, b, c, d, e, f, 2, 0xB5C0FBCF, v2) + ROUND_SECOND_BLOCK_CONST(f, g, h, a, b, c, d, e, 3, 0xE9B5DBA5, v3) 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) @@ -733,7 +624,6 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig 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) @@ -809,74 +699,9 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig schedule15 = first_hash7; ALL_SCHEDULE_LAST16() - - 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) - + ALL_ROUND_B1_1() ALL_SCHEDULE32() - - 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) + ALL_ROUND_B1_2() hash0 = UINT32_C(0x6A09E667) + a; hash1 = UINT32_C(0xBB67AE85) + b; @@ -896,10 +721,10 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig g = hash6; h = hash7; - 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_CONST(a, b, c, d, e, f, g, h, 0, 0x428A2F98, v0) + ROUND_SECOND_BLOCK_CONST(h, a, b, c, d, e, f, g, 1, 0x71374491, v1) + ROUND_SECOND_BLOCK_CONST(g, h, a, b, c, d, e, f, 2, 0xB5C0FBCF, v2) + ROUND_SECOND_BLOCK_CONST(f, g, h, a, b, c, d, e, 3, 0xE9B5DBA5, v3) 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) @@ -928,7 +753,6 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig 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) @@ -961,7 +785,7 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig 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) - + hash0 += a; hash1 += b; hash2 += c; @@ -974,10 +798,12 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig indexW += SINGLE_BLOCK_W_SIZE; } +//----------------------------------------------------- FINAL CHECK ------------------------------------------------ + schedule0 = __byte_perm(((uint32_t *)(IV))[0], 0, 0x0123) ^ hash0; - schedule1 = __byte_perm(((uint32_t *)(IV+4))[0], 0, 0x0123) ^ hash1; - schedule2 = __byte_perm(((uint32_t *)(IV+8))[0], 0, 0x0123) ^ hash2; - schedule3 = __byte_perm(((uint32_t *)(IV+12))[0], 0, 0x0123) ^ hash3; + schedule1 = __byte_perm(((uint32_t *)(IV+4))[0], 0, 0x0123) ^ hash1; + schedule2 = __byte_perm(((uint32_t *)(IV+8))[0], 0, 0x0123) ^ hash2; + schedule3 = __byte_perm(((uint32_t *)(IV+12))[0], 0, 0x0123) ^ hash3; schedule4 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule0 >> 24], TS1[(schedule1 >> 16) & 0xFF], TS2[(schedule2 >> 8) & 0xFF]) , TS3[schedule3 & 0xFF] , hash4); schedule5 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule1 >> 24], TS1[(schedule2 >> 16) & 0xFF], TS2[(schedule3 >> 8) & 0xFF]) , TS3[schedule0 & 0xFF] , hash5); @@ -1188,3 +1014,911 @@ __global__ void decrypt_vmk(int numStream, int tot_psw_kernel, int *found, unsig return; } + +__device__ void encrypt( + uint32_t k0, uint32_t k1, uint32_t k2, uint32_t k3, uint32_t k4, uint32_t k5, uint32_t k6, uint32_t k7, + uint32_t m0, uint32_t m1, uint32_t m2, uint32_t m3, + uint32_t * output0, uint32_t * output1, uint32_t * output2, uint32_t * output3 +) +{ + uint32_t enc_schedule0, enc_schedule1, enc_schedule2, enc_schedule3, enc_schedule4, enc_schedule5, enc_schedule6, enc_schedule7; + uint32_t local_key0, local_key1, local_key2, local_key3, local_key4, local_key5, local_key6, local_key7; + + local_key0=k0; + local_key1=k1; + local_key2=k2; + local_key3=k3; + local_key4=k4; + local_key5=k5; + local_key6=k6; + local_key7=k7; + + enc_schedule0 = __byte_perm(m0, 0, 0x0123) ^ local_key0; + enc_schedule1 = __byte_perm(m1, 0, 0x0123) ^ local_key1; + enc_schedule2 = __byte_perm(m2, 0, 0x0123) ^ local_key2; + enc_schedule3 = __byte_perm(m3, 0, 0x0123) ^ local_key3; + + enc_schedule4 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule0 >> 24], TS1[(enc_schedule1 >> 16) & 0xFF], TS2[(enc_schedule2 >> 8) & 0xFF]) , TS3[enc_schedule3 & 0xFF] , local_key4); + enc_schedule5 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule1 >> 24], TS1[(enc_schedule2 >> 16) & 0xFF], TS2[(enc_schedule3 >> 8) & 0xFF]) , TS3[enc_schedule0 & 0xFF] , local_key5); + enc_schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule2 >> 24], TS1[(enc_schedule3 >> 16) & 0xFF], TS2[(enc_schedule0 >> 8) & 0xFF]) , TS3[enc_schedule1 & 0xFF] , local_key6); + enc_schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule3 >> 24], TS1[(enc_schedule0 >> 16) & 0xFF], TS2[(enc_schedule1 >> 8) & 0xFF]) , TS3[enc_schedule2 & 0xFF] , local_key7); + + local_key0 ^= LOP3LUT_XOR( + LOP3LUT_XOR( (TS2[(local_key7 >> 24) ] & 0x000000FF), (TS3[(local_key7 >> 16) & 0xFF] & 0xFF000000), (TS0[(local_key7 >> 8) & 0xFF] & 0x00FF0000)), + (TS1[(local_key7 ) & 0xFF] & 0x0000FF00), 0x01000000 + ); //RCON[0]; + local_key1 ^= local_key0; local_key2 ^= local_key1; local_key3 ^= local_key2; + + enc_schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule4 >> 24], TS1[(enc_schedule5 >> 16) & 0xFF], TS2[(enc_schedule6 >> 8) & 0xFF]) , TS3[enc_schedule7 & 0xFF] , local_key0); + enc_schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule5 >> 24], TS1[(enc_schedule6 >> 16) & 0xFF], TS2[(enc_schedule7 >> 8) & 0xFF]) , TS3[enc_schedule4 & 0xFF] , local_key1); + enc_schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule6 >> 24], TS1[(enc_schedule7 >> 16) & 0xFF], TS2[(enc_schedule4 >> 8) & 0xFF]) , TS3[enc_schedule5 & 0xFF] , local_key2); + enc_schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule7 >> 24], TS1[(enc_schedule4 >> 16) & 0xFF], TS2[(enc_schedule5 >> 8) & 0xFF]) , TS3[enc_schedule6 & 0xFF] , local_key3); + + local_key4 ^= (TS3[(local_key3 >> 24) ] & 0xFF000000) ^ + (TS0[(local_key3 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key3 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS2[(local_key3 ) & 0xFF] & 0x000000FF); + local_key5 ^= local_key4; + local_key6 ^= local_key5; + local_key7 ^= local_key6; + + enc_schedule4 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule0 >> 24], TS1[(enc_schedule1 >> 16) & 0xFF], TS2[(enc_schedule2 >> 8) & 0xFF]) , TS3[enc_schedule3 & 0xFF] , local_key4); + enc_schedule5 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule1 >> 24], TS1[(enc_schedule2 >> 16) & 0xFF], TS2[(enc_schedule3 >> 8) & 0xFF]) , TS3[enc_schedule0 & 0xFF] , local_key5); + enc_schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule2 >> 24], TS1[(enc_schedule3 >> 16) & 0xFF], TS2[(enc_schedule0 >> 8) & 0xFF]) , TS3[enc_schedule1 & 0xFF] , local_key6); + enc_schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule3 >> 24], TS1[(enc_schedule0 >> 16) & 0xFF], TS2[(enc_schedule1 >> 8) & 0xFF]) , TS3[enc_schedule2 & 0xFF] , local_key7); + + local_key0 ^= (TS2[(local_key7 >> 24) ] & 0x000000FF) ^ + (TS3[(local_key7 >> 16) & 0xFF] & 0xFF000000) ^ + (TS0[(local_key7 >> 8) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key7 ) & 0xFF] & 0x0000FF00) ^ 0x02000000; //RCON[1]; + local_key1 ^= local_key0; local_key2 ^= local_key1; local_key3 ^= local_key2; + + enc_schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule4 >> 24], TS1[(enc_schedule5 >> 16) & 0xFF], TS2[(enc_schedule6 >> 8) & 0xFF]) , TS3[enc_schedule7 & 0xFF] , local_key0); + enc_schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule5 >> 24], TS1[(enc_schedule6 >> 16) & 0xFF], TS2[(enc_schedule7 >> 8) & 0xFF]) , TS3[enc_schedule4 & 0xFF] , local_key1); + enc_schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule6 >> 24], TS1[(enc_schedule7 >> 16) & 0xFF], TS2[(enc_schedule4 >> 8) & 0xFF]) , TS3[enc_schedule5 & 0xFF] , local_key2); + enc_schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule7 >> 24], TS1[(enc_schedule4 >> 16) & 0xFF], TS2[(enc_schedule5 >> 8) & 0xFF]) , TS3[enc_schedule6 & 0xFF] , local_key3); + + local_key4 ^= (TS3[(local_key3 >> 24) ] & 0xFF000000) ^ + (TS0[(local_key3 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key3 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS2[(local_key3 ) & 0xFF] & 0x000000FF); + local_key5 ^= local_key4; + local_key6 ^= local_key5; + local_key7 ^= local_key6; + + enc_schedule4 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule0 >> 24], TS1[(enc_schedule1 >> 16) & 0xFF], TS2[(enc_schedule2 >> 8) & 0xFF]) , TS3[enc_schedule3 & 0xFF] , local_key4); + enc_schedule5 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule1 >> 24], TS1[(enc_schedule2 >> 16) & 0xFF], TS2[(enc_schedule3 >> 8) & 0xFF]) , TS3[enc_schedule0 & 0xFF] , local_key5); + enc_schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule2 >> 24], TS1[(enc_schedule3 >> 16) & 0xFF], TS2[(enc_schedule0 >> 8) & 0xFF]) , TS3[enc_schedule1 & 0xFF] , local_key6); + enc_schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule3 >> 24], TS1[(enc_schedule0 >> 16) & 0xFF], TS2[(enc_schedule1 >> 8) & 0xFF]) , TS3[enc_schedule2 & 0xFF] , local_key7); + + + local_key0 ^= (TS2[(local_key7 >> 24) ] & 0x000000FF) ^ + (TS3[(local_key7 >> 16) & 0xFF] & 0xFF000000) ^ + (TS0[(local_key7 >> 8) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key7 ) & 0xFF] & 0x0000FF00) ^ 0x04000000; //RCON[2]; + local_key1 ^= local_key0; local_key2 ^= local_key1; local_key3 ^= local_key2; + + enc_schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule4 >> 24], TS1[(enc_schedule5 >> 16) & 0xFF], TS2[(enc_schedule6 >> 8) & 0xFF]) , TS3[enc_schedule7 & 0xFF] , local_key0); + enc_schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule5 >> 24], TS1[(enc_schedule6 >> 16) & 0xFF], TS2[(enc_schedule7 >> 8) & 0xFF]) , TS3[enc_schedule4 & 0xFF] , local_key1); + enc_schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule6 >> 24], TS1[(enc_schedule7 >> 16) & 0xFF], TS2[(enc_schedule4 >> 8) & 0xFF]) , TS3[enc_schedule5 & 0xFF] , local_key2); + enc_schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule7 >> 24], TS1[(enc_schedule4 >> 16) & 0xFF], TS2[(enc_schedule5 >> 8) & 0xFF]) , TS3[enc_schedule6 & 0xFF] , local_key3); + + + local_key4 ^= (TS3[(local_key3 >> 24) ] & 0xFF000000) ^ + (TS0[(local_key3 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key3 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS2[(local_key3 ) & 0xFF] & 0x000000FF); + local_key5 ^= local_key4; + local_key6 ^= local_key5; + local_key7 ^= local_key6; + + enc_schedule4 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule0 >> 24], TS1[(enc_schedule1 >> 16) & 0xFF], TS2[(enc_schedule2 >> 8) & 0xFF]) , TS3[enc_schedule3 & 0xFF] , local_key4); + enc_schedule5 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule1 >> 24], TS1[(enc_schedule2 >> 16) & 0xFF], TS2[(enc_schedule3 >> 8) & 0xFF]) , TS3[enc_schedule0 & 0xFF] , local_key5); + enc_schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule2 >> 24], TS1[(enc_schedule3 >> 16) & 0xFF], TS2[(enc_schedule0 >> 8) & 0xFF]) , TS3[enc_schedule1 & 0xFF] , local_key6); + enc_schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule3 >> 24], TS1[(enc_schedule0 >> 16) & 0xFF], TS2[(enc_schedule1 >> 8) & 0xFF]) , TS3[enc_schedule2 & 0xFF] , local_key7); + + local_key0 ^= (TS2[(local_key7 >> 24) ] & 0x000000FF) ^ + (TS3[(local_key7 >> 16) & 0xFF] & 0xFF000000) ^ + (TS0[(local_key7 >> 8) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key7 ) & 0xFF] & 0x0000FF00) ^ 0x08000000; //RCON[3]; + local_key1 ^= local_key0; local_key2 ^= local_key1; local_key3 ^= local_key2; + + enc_schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule4 >> 24], TS1[(enc_schedule5 >> 16) & 0xFF], TS2[(enc_schedule6 >> 8) & 0xFF]) , TS3[enc_schedule7 & 0xFF] , local_key0); + enc_schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule5 >> 24], TS1[(enc_schedule6 >> 16) & 0xFF], TS2[(enc_schedule7 >> 8) & 0xFF]) , TS3[enc_schedule4 & 0xFF] , local_key1); + enc_schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule6 >> 24], TS1[(enc_schedule7 >> 16) & 0xFF], TS2[(enc_schedule4 >> 8) & 0xFF]) , TS3[enc_schedule5 & 0xFF] , local_key2); + enc_schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule7 >> 24], TS1[(enc_schedule4 >> 16) & 0xFF], TS2[(enc_schedule5 >> 8) & 0xFF]) , TS3[enc_schedule6 & 0xFF] , local_key3); + + local_key4 ^= (TS3[(local_key3 >> 24) ] & 0xFF000000) ^ + (TS0[(local_key3 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key3 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS2[(local_key3 ) & 0xFF] & 0x000000FF); + local_key5 ^= local_key4; + local_key6 ^= local_key5; + local_key7 ^= local_key6; + + enc_schedule4 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule0 >> 24], TS1[(enc_schedule1 >> 16) & 0xFF], TS2[(enc_schedule2 >> 8) & 0xFF]) , TS3[enc_schedule3 & 0xFF] , local_key4); + enc_schedule5 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule1 >> 24], TS1[(enc_schedule2 >> 16) & 0xFF], TS2[(enc_schedule3 >> 8) & 0xFF]) , TS3[enc_schedule0 & 0xFF] , local_key5); + enc_schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule2 >> 24], TS1[(enc_schedule3 >> 16) & 0xFF], TS2[(enc_schedule0 >> 8) & 0xFF]) , TS3[enc_schedule1 & 0xFF] , local_key6); + enc_schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule3 >> 24], TS1[(enc_schedule0 >> 16) & 0xFF], TS2[(enc_schedule1 >> 8) & 0xFF]) , TS3[enc_schedule2 & 0xFF] , local_key7); + + local_key0 ^= (TS2[(local_key7 >> 24) ] & 0x000000FF) ^ + (TS3[(local_key7 >> 16) & 0xFF] & 0xFF000000) ^ + (TS0[(local_key7 >> 8) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key7 ) & 0xFF] & 0x0000FF00) ^ 0x10000000; //RCON[4]; + local_key1 ^= local_key0; local_key2 ^= local_key1; local_key3 ^= local_key2; + + enc_schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule4 >> 24], TS1[(enc_schedule5 >> 16) & 0xFF], TS2[(enc_schedule6 >> 8) & 0xFF]) , TS3[enc_schedule7 & 0xFF] , local_key0); + enc_schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule5 >> 24], TS1[(enc_schedule6 >> 16) & 0xFF], TS2[(enc_schedule7 >> 8) & 0xFF]) , TS3[enc_schedule4 & 0xFF] , local_key1); + enc_schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule6 >> 24], TS1[(enc_schedule7 >> 16) & 0xFF], TS2[(enc_schedule4 >> 8) & 0xFF]) , TS3[enc_schedule5 & 0xFF] , local_key2); + enc_schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule7 >> 24], TS1[(enc_schedule4 >> 16) & 0xFF], TS2[(enc_schedule5 >> 8) & 0xFF]) , TS3[enc_schedule6 & 0xFF] , local_key3); + + local_key4 ^= (TS3[(local_key3 >> 24) ] & 0xFF000000) ^ + (TS0[(local_key3 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key3 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS2[(local_key3 ) & 0xFF] & 0x000000FF); + local_key5 ^= local_key4; + local_key6 ^= local_key5; + local_key7 ^= local_key6; + + enc_schedule4 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule0 >> 24], TS1[(enc_schedule1 >> 16) & 0xFF], TS2[(enc_schedule2 >> 8) & 0xFF]) , TS3[enc_schedule3 & 0xFF] , local_key4); + enc_schedule5 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule1 >> 24], TS1[(enc_schedule2 >> 16) & 0xFF], TS2[(enc_schedule3 >> 8) & 0xFF]) , TS3[enc_schedule0 & 0xFF] , local_key5); + enc_schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule2 >> 24], TS1[(enc_schedule3 >> 16) & 0xFF], TS2[(enc_schedule0 >> 8) & 0xFF]) , TS3[enc_schedule1 & 0xFF] , local_key6); + enc_schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule3 >> 24], TS1[(enc_schedule0 >> 16) & 0xFF], TS2[(enc_schedule1 >> 8) & 0xFF]) , TS3[enc_schedule2 & 0xFF] , local_key7); + + + local_key0 ^= (TS2[(local_key7 >> 24) ] & 0x000000FF) ^ + (TS3[(local_key7 >> 16) & 0xFF] & 0xFF000000) ^ + (TS0[(local_key7 >> 8) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key7 ) & 0xFF] & 0x0000FF00) ^ 0x20000000; //RCON[5]; + local_key1 ^= local_key0; local_key2 ^= local_key1; local_key3 ^= local_key2; + + enc_schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule4 >> 24], TS1[(enc_schedule5 >> 16) & 0xFF], TS2[(enc_schedule6 >> 8) & 0xFF]) , TS3[enc_schedule7 & 0xFF] , local_key0); + enc_schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule5 >> 24], TS1[(enc_schedule6 >> 16) & 0xFF], TS2[(enc_schedule7 >> 8) & 0xFF]) , TS3[enc_schedule4 & 0xFF] , local_key1); + enc_schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule6 >> 24], TS1[(enc_schedule7 >> 16) & 0xFF], TS2[(enc_schedule4 >> 8) & 0xFF]) , TS3[enc_schedule5 & 0xFF] , local_key2); + enc_schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule7 >> 24], TS1[(enc_schedule4 >> 16) & 0xFF], TS2[(enc_schedule5 >> 8) & 0xFF]) , TS3[enc_schedule6 & 0xFF] , local_key3); + + local_key4 ^= (TS3[(local_key3 >> 24)] & 0xFF000000) ^ + (TS0[(local_key3 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key3 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS2[(local_key3 ) & 0xFF] & 0x000000FF); + local_key5 ^= local_key4; + local_key6 ^= local_key5; + local_key7 ^= local_key6; + + enc_schedule4 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule0 >> 24], TS1[(enc_schedule1 >> 16) & 0xFF], TS2[(enc_schedule2 >> 8) & 0xFF]) , TS3[enc_schedule3 & 0xFF] , local_key4); + enc_schedule5 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule1 >> 24], TS1[(enc_schedule2 >> 16) & 0xFF], TS2[(enc_schedule3 >> 8) & 0xFF]) , TS3[enc_schedule0 & 0xFF] , local_key5); + enc_schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule2 >> 24], TS1[(enc_schedule3 >> 16) & 0xFF], TS2[(enc_schedule0 >> 8) & 0xFF]) , TS3[enc_schedule1 & 0xFF] , local_key6); + enc_schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[enc_schedule3 >> 24], TS1[(enc_schedule0 >> 16) & 0xFF], TS2[(enc_schedule1 >> 8) & 0xFF]) , TS3[enc_schedule2 & 0xFF] , local_key7); + + local_key0 ^= (TS2[(local_key7 >> 24)] & 0x000000FF) ^ + (TS3[(local_key7 >> 16) & 0xFF] & 0xFF000000) ^ + (TS0[(local_key7 >> 8) & 0xFF] & 0x00FF0000) ^ + (TS1[(local_key7 ) & 0xFF] & 0x0000FF00) ^ 0x40000000; //RCON[6]; + local_key1 ^= local_key0; + local_key2 ^= local_key1; + local_key3 ^= local_key2; + + enc_schedule0 = (TS2[(enc_schedule4 >> 24) ] & 0xFF000000) ^ + (TS3[(enc_schedule5 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS0[(enc_schedule6 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS1[(enc_schedule7 ) & 0xFF] & 0x000000FF) ^ local_key0; + + enc_schedule1 = (TS2[(enc_schedule5 >> 24) ] & 0xFF000000) ^ + (TS3[(enc_schedule6 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS0[(enc_schedule7 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS1[(enc_schedule4 ) & 0xFF] & 0x000000FF) ^ local_key1; + + enc_schedule2 = (TS2[(enc_schedule6 >> 24) ] & 0xFF000000) ^ + (TS3[(enc_schedule7 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS0[(enc_schedule4 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS1[(enc_schedule5 ) & 0xFF] & 0x000000FF) ^ local_key2; + + enc_schedule3 = (TS2[(enc_schedule7 >> 24) ] & 0xFF000000) ^ + (TS3[(enc_schedule4 >> 16) & 0xFF] & 0x00FF0000) ^ + (TS0[(enc_schedule5 >> 8) & 0xFF] & 0x0000FF00) ^ + (TS1[(enc_schedule6 ) & 0xFF] & 0x000000FF) ^ local_key3; + + output0[0] = __byte_perm(enc_schedule0, 0, 0x0123); + output1[0] = __byte_perm(enc_schedule1, 0, 0x0123); + output2[0] = __byte_perm(enc_schedule2, 0, 0x0123); + output3[0] = __byte_perm(enc_schedule3, 0, 0x0123); + +} + +__global__ 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 +) +{ + uint32_t schedule0, schedule1, schedule2, schedule3, schedule4, schedule5, schedule6, schedule7, schedule8, schedule9; + uint32_t schedule10, schedule11, schedule12, schedule13, schedule14, schedule15, schedule16, schedule17, schedule18, schedule19; + uint32_t schedule20, schedule21, schedule22, schedule23, schedule24, schedule25, schedule26, schedule27, schedule28, schedule29; + uint32_t schedule30, schedule31; + uint32_t first_hash0, first_hash1, first_hash2, first_hash3, first_hash4, first_hash5, first_hash6, first_hash7; + uint32_t hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7; + uint32_t a, b, c, d, e, f, g, h; + + int gIndex = (threadIdx.x+blockIdx.x*blockDim.x); + int index_generic; + int indexW=(gIndex*FIXED_PASSWORD_BUFFER); + int curr_fetch=0; + + while(gIndex < tot_psw_kernel) + { + + first_hash0 = UINT32_C(0x6A09E667); + first_hash1 = UINT32_C(0xBB67AE85); + first_hash2 = UINT32_C(0x3C6EF372); + first_hash3 = UINT32_C(0xA54FF53A); + first_hash4 = UINT32_C(0x510E527F); + first_hash5 = UINT32_C(0x9B05688C); + first_hash6 = UINT32_C(0x1F83D9AB); + first_hash7 = UINT32_C(0x5BE0CD19); + + a = UINT32_C(0x6A09E667); + b = UINT32_C(0xBB67AE85); + c = UINT32_C(0x3C6EF372); + d = UINT32_C(0xA54FF53A); + e = UINT32_C(0x510E527F); + f = UINT32_C(0x9B05688C); + g = UINT32_C(0x1F83D9AB); + h = UINT32_C(0x5BE0CD19); + +//----------------------------------------------------- FIRST HASH ------------------------------------------------ + indexW=(gIndex*FIXED_PASSWORD_BUFFER); + curr_fetch=0; + index_generic=MAX_INPUT_PASSWORD_LEN; + if(numStream == 0) + { + schedule0 = ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch+1)) << 8) | 0; + curr_fetch+=2; + schedule1 = ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch+1)) << 8) | 0; + curr_fetch+=2; + schedule2 = ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch+1)) << 8) | 0; + curr_fetch+=2; + schedule3 = ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch+1)) << 8) | 0; + curr_fetch+=2; + + schedule4 = ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password0, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password0, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule5 = ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password0, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch;} + if(tex1Dfetch(w_password0, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1;} + curr_fetch+=2; + + schedule6 = ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password0, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password0, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule7 = ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password0, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password0, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule8 = ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password0, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password0, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule9 = ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password0, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password0, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule10 = ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password0, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password0, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule11 = ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password0, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password0, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule12 = ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password0, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password0, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule13 = ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password0, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password0, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password0, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + } + else + { + schedule0 = ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch+1)) << 8) | 0; + curr_fetch+=2; + schedule1 = ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch+1)) << 8) | 0; + curr_fetch+=2; + schedule2 = ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch+1)) << 8) | 0; + curr_fetch+=2; + schedule3 = ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch+1)) << 8) | 0; + curr_fetch+=2; + + schedule4 = ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password1, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password1, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule5 = ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password1, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password1, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule6 = ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password1, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password1, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule7 = ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password1, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password1, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule8 = ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password1, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password1, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule9 = ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password1, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password1, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule10 = ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password1, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password1, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule11 = ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password1, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password1, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule12 = ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password1, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password1, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + curr_fetch+=2; + + schedule13 = ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch)) << 24) | 0 | ((uint32_t)tex1Dfetch(w_password1, (indexW+curr_fetch+1)) << 8) | 0; + if(tex1Dfetch(w_password1, (indexW+curr_fetch)) == END_STRING) { index_generic=curr_fetch; } + if(tex1Dfetch(w_password1, (indexW+curr_fetch+1)) == END_STRING) { index_generic=curr_fetch+1; } + } + + + if(index_generic == MAX_INPUT_PASSWORD_LEN) schedule13 = schedule13 | ((uint32_t)0x8000); + + schedule14=0; + index_generic*=2; + schedule15 = ((uint8_t)((index_generic << 3) >> 8)) << 8 | ((uint8_t)(index_generic << 3)); + + ALL_SCHEDULE_LAST16() + ALL_ROUND_B1_1() + ALL_SCHEDULE32() + ALL_ROUND_B1_2() + + first_hash0 += a; + first_hash1 += b; + first_hash2 += c; + first_hash3 += d; + first_hash4 += e; + first_hash5 += f; + first_hash6 += g; + first_hash7 += h; + +//----------------------------------------------------- SECOND HASH ------------------------------------------------ + //old loadschedule + schedule0 = first_hash0; + schedule1 = first_hash1; + schedule2 = first_hash2; + schedule3 = first_hash3; + schedule4 = first_hash4; + schedule5 = first_hash5; + schedule6 = first_hash6; + schedule7 = first_hash7; + schedule8 = 0x80000000; + schedule9 = 0; + schedule10 = 0; + schedule11 = 0; + schedule12 = 0; + schedule13 = 0; + schedule14 = 0; + schedule15 = 0x100; + + first_hash0 = UINT32_C(0x6A09E667); + first_hash1 = UINT32_C(0xBB67AE85); + first_hash2 = UINT32_C(0x3C6EF372); + first_hash3 = UINT32_C(0xA54FF53A); + first_hash4 = UINT32_C(0x510E527F); + first_hash5 = UINT32_C(0x9B05688C); + first_hash6 = UINT32_C(0x1F83D9AB); + first_hash7 = UINT32_C(0x5BE0CD19); + + a = first_hash0; + b = first_hash1; + c = first_hash2; + d = first_hash3; + e = first_hash4; + f = first_hash5; + g = first_hash6; + h = first_hash7; + + ALL_SCHEDULE_LAST16() + ALL_ROUND_B1_1() + ALL_SCHEDULE32() + ALL_ROUND_B1_2() + + first_hash0 += a; + first_hash1 += b; + first_hash2 += c; + first_hash3 += d; + first_hash4 += e; + first_hash5 += f; + first_hash6 += g; + first_hash7 += h; + +//----------------------------------------------------- LOOP HASH ------------------------------------------------ + + hash0=0; + hash1=0; + hash2=0; + hash3=0; + hash4=0; + hash5=0; + hash6=0; + hash7=0; + + indexW=0; + + for(index_generic=0; index_generic < ITERATION_NUMBER/2; index_generic++) + { + a = UINT32_C(0x6A09E667); + b = UINT32_C(0xBB67AE85); + c = UINT32_C(0x3C6EF372); + d = UINT32_C(0xA54FF53A); + e = UINT32_C(0x510E527F); + f = UINT32_C(0x9B05688C); + g = UINT32_C(0x1F83D9AB); + h = UINT32_C(0x5BE0CD19); + + schedule0 = hash0; + schedule1 = hash1; + schedule2 = hash2; + schedule3 = hash3; + schedule4 = hash4; + schedule5 = hash5; + schedule6 = hash6; + schedule7 = hash7; + + schedule8 = first_hash0; + schedule9 = first_hash1; + schedule10 = first_hash2; + schedule11 = first_hash3; + schedule12 = first_hash4; + schedule13 = first_hash5; + schedule14 = first_hash6; + schedule15 = first_hash7; + + ALL_SCHEDULE_LAST16() + ALL_ROUND_B1_1() + ALL_SCHEDULE32() + ALL_ROUND_B1_2() + + hash0 = UINT32_C(0x6A09E667) + a; + hash1 = UINT32_C(0xBB67AE85) + b; + hash2 = UINT32_C(0x3C6EF372) + c; + hash3 = UINT32_C(0xA54FF53A) + d; + hash4 = UINT32_C(0x510E527F) + e; + hash5 = UINT32_C(0x9B05688C) + f; + hash6 = UINT32_C(0x1F83D9AB) + g; + hash7 = UINT32_C(0x5BE0CD19) + h; + + a = hash0; + b = hash1; + c = hash2; + d = hash3; + e = hash4; + f = hash5; + g = hash6; + h = hash7; + + ROUND_SECOND_BLOCK_CONST(a, b, c, d, e, f, g, h, 0, 0x428A2F98, v0) + ROUND_SECOND_BLOCK_CONST(h, a, b, c, d, e, f, g, 1, 0x71374491, v1) + ROUND_SECOND_BLOCK_CONST(g, h, a, b, c, d, e, f, 2, 0xB5C0FBCF, v2) + ROUND_SECOND_BLOCK_CONST(f, g, h, a, b, c, d, e, 3, 0xE9B5DBA5, v3) + + 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) + + hash0 += a; + hash1 += b; + hash2 += c; + hash3 += d; + hash4 += e; + hash5 += f; + hash6 += g; + hash7 += h; + + indexW += SINGLE_BLOCK_W_SIZE; + } + + for(index_generic=ITERATION_NUMBER/2; index_generic < ITERATION_NUMBER; index_generic++) + { + a = UINT32_C(0x6A09E667); + b = UINT32_C(0xBB67AE85); + c = UINT32_C(0x3C6EF372); + d = UINT32_C(0xA54FF53A); + e = UINT32_C(0x510E527F); + f = UINT32_C(0x9B05688C); + g = UINT32_C(0x1F83D9AB); + h = UINT32_C(0x5BE0CD19); + + schedule0 = hash0; + schedule1 = hash1; + schedule2 = hash2; + schedule3 = hash3; + schedule4 = hash4; + schedule5 = hash5; + schedule6 = hash6; + schedule7 = hash7; + + schedule8 = first_hash0; + schedule9 = first_hash1; + schedule10 = first_hash2; + schedule11 = first_hash3; + schedule12 = first_hash4; + schedule13 = first_hash5; + schedule14 = first_hash6; + schedule15 = first_hash7; + + ALL_SCHEDULE_LAST16() + ALL_ROUND_B1_1() + ALL_SCHEDULE32() + ALL_ROUND_B1_2() + + hash0 = UINT32_C(0x6A09E667) + a; + hash1 = UINT32_C(0xBB67AE85) + b; + hash2 = UINT32_C(0x3C6EF372) + c; + hash3 = UINT32_C(0xA54FF53A) + d; + hash4 = UINT32_C(0x510E527F) + e; + hash5 = UINT32_C(0x9B05688C) + f; + hash6 = UINT32_C(0x1F83D9AB) + g; + hash7 = UINT32_C(0x5BE0CD19) + h; + + a = hash0; + b = hash1; + c = hash2; + d = hash3; + e = hash4; + f = hash5; + g = hash6; + h = hash7; + + ROUND_SECOND_BLOCK_CONST(a, b, c, d, e, f, g, h, 0, 0x428A2F98, v0) + ROUND_SECOND_BLOCK_CONST(h, a, b, c, d, e, f, g, 1, 0x71374491, v1) + ROUND_SECOND_BLOCK_CONST(g, h, a, b, c, d, e, f, 2, 0xB5C0FBCF, v2) + ROUND_SECOND_BLOCK_CONST(f, g, h, a, b, c, d, e, 3, 0xE9B5DBA5, v3) + + 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) + + hash0 += a; + hash1 += b; + hash2 += c; + hash3 += d; + hash4 += e; + hash5 += f; + hash6 += g; + hash7 += h; + + indexW += SINGLE_BLOCK_W_SIZE; + } + +//----------------------------------------------------- MAC COMPARISON ------------------------------------------------ + + a = ((uint32_t *)(vmkIV))[0]; + b = ((uint32_t *)(vmkIV+4))[0]; + c = ((uint32_t *)(vmkIV+8))[0]; + d = ((uint32_t *)(vmkIV+12))[0]; + + encrypt( + hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7, + a, b, c, d, + &(schedule0), &(schedule1), &(schedule2), &(schedule3) + ); + + schedule0= + (((uint32_t)(vmkKey[3] ^ ((uint8_t) (schedule0 >> 24) ))) << 24) | + (((uint32_t)(vmkKey[2] ^ ((uint8_t) (schedule0 >> 16) ))) << 16) | + (((uint32_t)(vmkKey[1] ^ ((uint8_t) (schedule0 >> 8) ))) << 8) | + (((uint32_t)(vmkKey[0] ^ ((uint8_t) (schedule0)))) << 0); + + schedule1= + (((uint32_t)(vmkKey[7] ^ ((uint8_t) (schedule1 >> 24) ))) << 24) | + (((uint32_t)(vmkKey[6] ^ ((uint8_t) (schedule1 >> 16) ))) << 16) | + (((uint32_t)(vmkKey[5] ^ ((uint8_t) (schedule1 >> 8) ))) << 8) | + (((uint32_t)(vmkKey[4] ^ ((uint8_t) (schedule1)))) << 0); + + schedule2= + (((uint32_t)(vmkKey[11] ^ ((uint8_t) (schedule2 >> 24) ))) << 24) | + (((uint32_t)(vmkKey[10] ^ ((uint8_t) (schedule2 >> 16) ))) << 16) | + (((uint32_t)(vmkKey[9] ^ ((uint8_t) (schedule2 >> 8) ))) << 8) | + (((uint32_t)(vmkKey[8] ^ ((uint8_t) (schedule2)))) << 0); + + schedule3= + (((uint32_t)(vmkKey[15] ^ ((uint8_t) (schedule3 >> 24) ))) << 24) | + (((uint32_t)(vmkKey[14] ^ ((uint8_t) (schedule3 >> 16) ))) << 16) | + (((uint32_t)(vmkKey[13] ^ ((uint8_t) (schedule3 >> 8) ))) << 8) | + (((uint32_t)(vmkKey[12] ^ ((uint8_t) (schedule3)))) << 0); + + d += 0x01000000; + + encrypt( + hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7, + a, b, c, d, + &(schedule4), &(schedule5), &(schedule6), &(schedule7) + ); + + schedule4= + (((uint32_t)(vmkKey[19] ^ ((uint8_t) (schedule4 >> 24) ))) << 24) | + (((uint32_t)(vmkKey[18] ^ ((uint8_t) (schedule4 >> 16) ))) << 16) | + (((uint32_t)(vmkKey[17] ^ ((uint8_t) (schedule4 >> 8) ))) << 8) | + (((uint32_t)(vmkKey[16] ^ ((uint8_t) (schedule4)))) << 0); + + schedule5= + (((uint32_t)(vmkKey[23] ^ ((uint8_t) (schedule5 >> 24) ))) << 24) | + (((uint32_t)(vmkKey[22] ^ ((uint8_t) (schedule5 >> 16) ))) << 16) | + (((uint32_t)(vmkKey[21] ^ ((uint8_t) (schedule5 >> 8) ))) << 8) | + (((uint32_t)(vmkKey[20] ^ ((uint8_t) (schedule5)))) << 0); + + schedule6= + (((uint32_t)(vmkKey[27] ^ ((uint8_t) (schedule6 >> 24) ))) << 24) | + (((uint32_t)(vmkKey[26] ^ ((uint8_t) (schedule6 >> 16) ))) << 16) | + (((uint32_t)(vmkKey[25] ^ ((uint8_t) (schedule6 >> 8) ))) << 8) | + (((uint32_t)(vmkKey[24] ^ ((uint8_t) (schedule6)))) << 0); + + schedule7= + (((uint32_t)(vmkKey[31] ^ ((uint8_t) (schedule7 >> 24) ))) << 24) | + (((uint32_t)(vmkKey[30] ^ ((uint8_t) (schedule7 >> 16) ))) << 16) | + (((uint32_t)(vmkKey[29] ^ ((uint8_t) (schedule7 >> 8) ))) << 8) | + (((uint32_t)(vmkKey[28] ^ ((uint8_t) (schedule7)))) << 0); + + d += 0x01000000; + + encrypt( + hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7, + a, b, c, d, + &(schedule8), &(schedule9), &(schedule10), &(schedule11) + ); + + schedule8= + (((uint32_t)(vmkKey[35] ^ ((uint8_t) (schedule8 >> 24) ))) << 24) | + (((uint32_t)(vmkKey[34] ^ ((uint8_t) (schedule8 >> 16) ))) << 16) | + (((uint32_t)(vmkKey[33] ^ ((uint8_t) (schedule8 >> 8) ))) << 8) | + (((uint32_t)(vmkKey[32] ^ ((uint8_t) (schedule8)))) << 0); + + schedule9= + (((uint32_t)(vmkKey[39] ^ ((uint8_t) (schedule9 >> 24) ))) << 24) | + (((uint32_t)(vmkKey[38] ^ ((uint8_t) (schedule9 >> 16) ))) << 16) | + (((uint32_t)(vmkKey[37] ^ ((uint8_t) (schedule9 >> 8) ))) << 8) | + (((uint32_t)(vmkKey[36] ^ ((uint8_t) (schedule9)))) << 0); + + schedule10= + (((uint32_t)(vmkKey[43] ^ ((uint8_t) (schedule10 >> 24) ))) << 24) | + (((uint32_t)(vmkKey[42] ^ ((uint8_t) (schedule10 >> 16) ))) << 16) | + (((uint32_t)(vmkKey[41] ^ ((uint8_t) (schedule10 >> 8) ))) << 8) | + (((uint32_t)(vmkKey[40] ^ ((uint8_t) (schedule10)))) << 0); + + encrypt( + hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7, + ((uint32_t *)(macIV))[0], ((uint32_t *)(macIV+4))[0], ((uint32_t *)(macIV+8))[0], ((uint32_t *)(macIV+12))[0], + &(schedule16), &(schedule17), &(schedule18), &(schedule19) + ); + + + encrypt( + hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7, + ((uint32_t *)(computeMacIV))[0], ((uint32_t *)(computeMacIV+4))[0], ((uint32_t *)(computeMacIV+8))[0], ((uint32_t *)(computeMacIV+12))[0], + &(schedule12), &(schedule13), &(schedule14), &(schedule15) + ); + + schedule28 = schedule0 ^ schedule12; + schedule29 = schedule1 ^ schedule13; + schedule30 = schedule2 ^ schedule14; + schedule31 = schedule3 ^ schedule15; + + encrypt( + hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7, + schedule28, schedule29, schedule30, schedule31, + &(schedule12), &(schedule13), &(schedule14), &(schedule15) + ); + + schedule28 = schedule4 ^ schedule12; + schedule29 = schedule5 ^ schedule13; + schedule30 = schedule6 ^ schedule14; + schedule31 = schedule7 ^ schedule15; + + encrypt( + hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7, + schedule28, schedule29, schedule30, schedule31, + &(schedule12), &(schedule13), &(schedule14), &(schedule15) + ); + + schedule28 = schedule8 ^ schedule12; + schedule29 = schedule9 ^ schedule13; + schedule30 = schedule10 ^ schedule14; + schedule31 = schedule15; + + encrypt( + hash0, hash1, hash2, hash3, hash4, hash5, hash6, hash7, + schedule28, schedule29, schedule30, schedule31, + &(schedule12), &(schedule13), &(schedule14), &(schedule15) + ); + + if ( + + ( + schedule12 == ( (uint32_t) + (((uint32_t)(mac[3] ^ ((uint8_t) (schedule16 >> 24) ))) << 24) | + (((uint32_t)(mac[2] ^ ((uint8_t) (schedule16 >> 16) ))) << 16) | + (((uint32_t)(mac[1] ^ ((uint8_t) (schedule16 >> 8) ))) << 8) | + (((uint32_t)(mac[0] ^ ((uint8_t) (schedule16)))) << 0) ) + ) + && + ( + schedule13 == ( (uint32_t) + (((uint32_t)(mac[7] ^ ((uint8_t) (schedule17 >> 24) ))) << 24) | + (((uint32_t)(mac[6] ^ ((uint8_t) (schedule17 >> 16) ))) << 16) | + (((uint32_t)(mac[5] ^ ((uint8_t) (schedule17 >> 8) ))) << 8) | + (((uint32_t)(mac[4] ^ ((uint8_t) (schedule17)))) << 0) ) + ) + && + ( + schedule14 == ( (uint32_t) + (((uint32_t)(mac[11] ^ ((uint8_t) (schedule18 >> 24) ))) << 24) | + (((uint32_t)(mac[10] ^ ((uint8_t) (schedule18 >> 16) ))) << 16) | + (((uint32_t)(mac[9] ^ ((uint8_t) (schedule18 >> 8) ))) << 8) | + (((uint32_t)(mac[8] ^ ((uint8_t) (schedule18)))) << 0) ) + ) + && + ( + schedule15 == ( (uint32_t) + (((uint32_t)(mac[15] ^ ((uint8_t) (schedule19 >> 24) ))) << 24) | + (((uint32_t)(mac[14] ^ ((uint8_t) (schedule19 >> 16) ))) << 16) | + (((uint32_t)(mac[13] ^ ((uint8_t) (schedule19 >> 8) ))) << 8) | + (((uint32_t)(mac[12] ^ ((uint8_t) (schedule19)))) << 0) ) + ) + ) + { + *found = gIndex; + break; + } + + gIndex += (blockDim.x * gridDim.x); + } + + return; +} diff --git a/CUDA_version/main.cu b/CUDA_version/main.cu index 0c4ee3d..3cd532e 100755 --- a/CUDA_version/main.cu +++ b/CUDA_version/main.cu @@ -3,7 +3,7 @@ * Copyright (C) 2013-2017 Elena Ago * Massimo Bernaschi * - * 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 -d \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: diff --git a/CUDA_version/sha256_header.h b/CUDA_version/sha256_header.h index 8a29be2..7648f88 100755 --- a/CUDA_version/sha256_header.h +++ b/CUDA_version/sha256_header.h @@ -3,7 +3,7 @@ * Copyright (C) 2013-2017 Elena Ago * Massimo Bernaschi * - * 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) diff --git a/CUDA_version/utils.cu b/CUDA_version/utils.cu index 66d42dd..4071396 100755 --- a/CUDA_version/utils.cu +++ b/CUDA_version/utils.cu @@ -3,7 +3,7 @@ * Copyright (C) 2013-2017 Elena Ago * Massimo Bernaschi * - * 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; } diff --git a/CUDA_version/w_blocks.cu b/CUDA_version/w_blocks.cu index f89ae9f..ad766ef 100755 --- a/CUDA_version/w_blocks.cu +++ b/CUDA_version/w_blocks.cu @@ -3,7 +3,7 @@ * Copyright (C) 2013-2017 Elena Ago * Massimo Bernaschi * - * 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); } } diff --git a/Hash_Extractor/Makefile b/Hash_Extractor/Makefile index 5472c3a..44137f1 100755 --- a/Hash_Extractor/Makefile +++ b/Hash_Extractor/Makefile @@ -3,4 +3,4 @@ bitcracker_hash: clean: rm -rf *.o - rm -rf bitcracker_hash + rm -rf ../build/bitcracker_hash diff --git a/Hash_Extractor/bitcracker_hash.c b/Hash_Extractor/bitcracker_hash.c index d426ce1..47551dc 100644 --- a/Hash_Extractor/bitcracker_hash.c +++ b/Hash_Extractor/bitcracker_hash.c @@ -70,12 +70,12 @@ void * Calloc(size_t len, size_t size) { static int usage(char *name){ - printf("\nUsage: %s -i -o \n\n" + printf("\nUsage: %s -i -o \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"); diff --git a/OpenCL_version/bitcracker.h b/OpenCL_version/bitcracker.h index 71ff2a0..ba49ea9 100755 --- a/OpenCL_version/bitcracker.h +++ b/OpenCL_version/bitcracker.h @@ -3,7 +3,7 @@ * Copyright (C) 2013-2017 Elena Ago * Massimo Bernaschi * - * 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 diff --git a/OpenCL_version/kernel_attack.cl b/OpenCL_version/kernel_attack.cl index 9182942..a99340a 100755 --- a/OpenCL_version/kernel_attack.cl +++ b/OpenCL_version/kernel_attack.cl @@ -3,7 +3,7 @@ * Copyright (C) 2013-2017 Elena Ago * Massimo Bernaschi * - * 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 diff --git a/OpenCL_version/kernel_wblocks.cl b/OpenCL_version/kernel_wblocks.cl index 44ad376..3f50528 100755 --- a/OpenCL_version/kernel_wblocks.cl +++ b/OpenCL_version/kernel_wblocks.cl @@ -3,7 +3,7 @@ * Copyright (C) 2013-2017 Elena Ago * Massimo Bernaschi * - * 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 diff --git a/OpenCL_version/main.c b/OpenCL_version/main.c index 05cf2de..511bc55 100755 --- a/OpenCL_version/main.c +++ b/OpenCL_version/main.c @@ -3,7 +3,7 @@ * Copyright (C) 2013-2017 Elena Ago * Massimo Bernaschi * - * 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 diff --git a/OpenCL_version/opencl_attack.c b/OpenCL_version/opencl_attack.c index b9f7275..bc35a75 100755 --- a/OpenCL_version/opencl_attack.c +++ b/OpenCL_version/opencl_attack.c @@ -3,7 +3,7 @@ * Copyright (C) 2013-2017 Elena Ago * Massimo Bernaschi * - * 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 diff --git a/OpenCL_version/utils.c b/OpenCL_version/utils.c index 42c0868..1565be2 100755 --- a/OpenCL_version/utils.c +++ b/OpenCL_version/utils.c @@ -3,7 +3,7 @@ * Copyright (C) 2013-2017 Elena Ago * Massimo Bernaschi * - * 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