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

1752 lines
80 KiB
Plaintext
Executable File
Raw Permalink Blame History

/*
* BitCracker: BitLocker password cracking tool, CUDA version.
* Copyright (C) 2013-2017 Elena Ago <elena dot ago at gmail dot com>
* Massimo Bernaschi <massimo dot bernaschi at gmail dot com>
*
* This file is part of the BitCracker project: https://github.com/e-ago/bitcracker
*
* BitCracker is free software: you can redistribute it and/or modify
* it under the terms of the GNU General Public License as published by
* the Free Software Foundation, either version 2 of the License, or
* (at your option) any later version.
*
* BitCracker is distributed in the hope that it will be useful,
* but WITHOUT ANY WARRANTY; without even the implied warranty of
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
* GNU General Public License for more details.
*
* You should have received a copy of the GNU General Public License
* along with BitCracker. If not, see <http://www.gnu.org/licenses/>.
*/
#include "bitcracker.h"
texture<uint32_t> w_texture;
texture<uint32_t> w_password;
int *deviceFound, *hostFound;
char *hostPassword;
uint32_t *hostPasswordInt, *devicePasswordInt;
unsigned char outPsw[MAX_INPUT_PASSWORD_LEN+1];
int outIndexPsw=0;
static int check_match() {
int i=0;
if (hostFound[0] >= 0){
outIndexPsw=(hostFound[0]);
snprintf((char*)outPsw, PSW_CHAR_SIZE, "%s", (char *)(hostPassword+(outIndexPsw*PSW_CHAR_SIZE)));
for(i=0; i<MAX_INPUT_PASSWORD_LEN; i++)
if(outPsw[i] == 0x80 || outPsw[i] == 0xffffff80) outPsw[i]='\0';
return 1;
}
return 0;
}
char *cuda_attack(
char *dname, uint32_t * w_blocks_d,
unsigned char * encryptedVMK,
unsigned char * nonce, unsigned char * encryptedMAC,
int gridBlocks)
{
FILE *fp;
int numReadPassword, match=0, done=0, w_blocks_h[4], cudaThreads=CUDA_THREADS_NO_MAC;
long long totReadPsw = 0;
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;
cudaEvent_t start, stop;
cudaStream_t stream;
float elapsedTime;
if(dname == NULL || w_blocks_d == NULL || encryptedVMK == NULL)
{
fprintf(stderr, "Attack input error\n");
return NULL;
}
if(tot_psw <= 0)
{
fprintf(stderr, "Attack tot passwords error: %d\n", tot_psw);
return NULL;
}
//-------- vmkIV setup ------
memset(vmkIV, 0, IV_SIZE);
vmkIV[0] = (unsigned char)(IV_SIZE - 1 - NONCE_SIZE - 1);
memcpy(vmkIV + 1, nonce, NONCE_SIZE);
if(IV_SIZE-1 - NONCE_SIZE - 1 < 0)
{
fprintf(stderr, "Attack nonce error\n");
return NULL;
}
vmkIV[IV_SIZE-1] = 1;
// -----------------------
if(mac_comparison == 1)
{
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;
} else {
fp = fopen(dname, "r");
if (!fp) {
fprintf(stderr, "Can't open dictionary file %s.\n", dname);
return NULL;
}
}
// -------------------------------
// ---- HOST VARS ----
BITCRACKER_CUDA_CHECK( cudaHostAlloc( (void ** ) &hostPasswordInt, tot_psw * PSW_INT_SIZE * sizeof(uint32_t), cudaHostAllocDefault) );
memset(hostPasswordInt, tot_psw * PSW_INT_SIZE * sizeof(uint32_t), 0);
BITCRACKER_CUDA_CHECK( cudaHostAlloc( (void ** ) &hostPassword, tot_psw*PSW_CHAR_SIZE*sizeof(char), cudaHostAllocDefault) );
BITCRACKER_CUDA_CHECK( cudaHostAlloc( (void ** ) &hostFound, sizeof(uint32_t), cudaHostAllocDefault) );
*hostFound = -1;
// ------------------------
// ---- CUDA VARS ----
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) );
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) );
}
BITCRACKER_CUDA_CHECK( cudaMalloc( (void ** ) &devicePasswordInt, (tot_psw * PSW_INT_SIZE * sizeof(uint32_t)) ) );
BITCRACKER_CUDA_CHECK( cudaMalloc( (void ** ) &deviceFound, (sizeof(uint32_t)) ) );
BITCRACKER_CUDA_CHECK( cudaMemcpy(deviceFound, hostFound, sizeof(uint32_t), cudaMemcpyHostToDevice) );
BITCRACKER_CUDA_CHECK( cudaStreamCreate(&(stream)) );
BITCRACKER_CUDA_CHECK( cudaEventCreate(&start) );
BITCRACKER_CUDA_CHECK( cudaEventCreate(&stop) );
// ---------------------
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_password, devicePasswordInt, (tot_psw * PSW_INT_SIZE * sizeof(uint32_t))));
// -------------------------
BITCRACKER_CUDA_CHECK (cudaDeviceSetCacheConfig( cudaFuncCachePreferL1 ) );
printf("Type of attack: %s\nCUDA Threads: %d\nCUDA Blocks: %d\nPsw per thread: %d\nMax Psw per kernel: %d\nDictionary: %s\nStrict Check (-s): %s\nMAC Comparison (-m): %s\n\n",
(attack_mode==MODE_USER_PASS)?"User Password":"Recovery Password", cudaThreads, gridBlocks, psw_x_thread, tot_psw, (fp == stdin)?"standard input":dname, (strict_check == 1)?"Yes":"No", (mac_comparison == 1)?"Yes":"No");
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]);
while(!done) {
numReadPassword = readFilePassword(&hostPasswordInt, &hostPassword, tot_psw, fp);
if(numReadPassword <= 0) { done=1; continue; }
BITCRACKER_CUDA_CHECK( cudaMemcpyAsync(devicePasswordInt, hostPasswordInt, tot_psw * PSW_INT_SIZE * sizeof(uint32_t), cudaMemcpyHostToDevice, stream) );
BITCRACKER_CUDA_CHECK( cudaEventRecord(start, stream) );
if(mac_comparison == 1)
{
//Slower attack with MAC verification
decrypt_vmk_with_mac<<<gridBlocks, CUDA_THREADS_WITH_MAC, 0, stream>>>(
numReadPassword, deviceFound,
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, attack_mode
);
}
else
{
//Faster attack
decrypt_vmk<<<gridBlocks, CUDA_THREADS_NO_MAC, 0, stream>>>(
numReadPassword, deviceFound, 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, attack_mode);
}
BITCRACKER_CUDA_CHECK_LAST_ERROR();
BITCRACKER_CUDA_CHECK( cudaEventRecord(stop, stream) );
BITCRACKER_CUDA_CHECK( cudaMemcpyAsync(hostFound, deviceFound, sizeof(unsigned int), cudaMemcpyDeviceToHost, stream) );
BITCRACKER_CUDA_CHECK( cudaStreamSynchronize(stream) );
totReadPsw += numReadPassword;
BITCRACKER_CUDA_CHECK( cudaEventElapsedTime(&elapsedTime, start, stop) );
printf("CUDA Kernel execution:\n\tEffective passwords: %d\n\tPasswords Range:\n\t\t%s\n\t\t.....\n\t\t%s\n\tTime: %f sec\n\tPasswords x second: %8.2f pw/sec\n",
numReadPassword,
(char *)hostPassword,
(char *)(hostPassword+((numReadPassword-1)*PSW_CHAR_SIZE)),
(elapsedTime/1000.0), numReadPassword/(elapsedTime/1000.0));
match=check_match();
if(match) done=1;
if(feof(fp)) done=1;
}
if (fp != stdin)
fclose(fp);
if(match==1)
printf("\n\n================================================\nCUDA attack completed\nPasswords evaluated: %lld\nPassword found: %s\n================================================\n\n", totReadPsw, outPsw);
else
printf("\n\n================================================\nCUDA attack completed\nPasswords evaluated: %lld\nPassword not found!\n================================================\n\n", totReadPsw);
BITCRACKER_CUDA_CHECK( cudaUnbindTexture(&w_password) );
BITCRACKER_CUDA_CHECK( cudaFreeHost(hostPassword) );
BITCRACKER_CUDA_CHECK( cudaFree(devicePasswordInt) );
BITCRACKER_CUDA_CHECK( cudaFree(deviceFound) );
BITCRACKER_CUDA_CHECK( cudaStreamDestroy(stream) );
BITCRACKER_CUDA_CHECK( cudaUnbindTexture(&w_texture) );
return NULL;
}
#define END_STRING 0x80 //0xFF
__global__ void decrypt_vmk(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, int method
)
{
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*PSW_INT_SIZE);
int8_t redo=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*PSW_INT_SIZE);
redo=0;
schedule0 = (uint32_t) (tex1Dfetch(w_password, (indexW+0)));
schedule1 = (uint32_t) (tex1Dfetch(w_password, (indexW+1)));
schedule2 = (uint32_t) (tex1Dfetch(w_password, (indexW+2)));
schedule3 = (uint32_t) (tex1Dfetch(w_password, (indexW+3)));
schedule4 = (uint32_t) (tex1Dfetch(w_password, (indexW+4)));
schedule5 = (uint32_t) (tex1Dfetch(w_password, (indexW+5)));
schedule6 = (uint32_t) (tex1Dfetch(w_password, (indexW+6)));
schedule7 = (uint32_t) (tex1Dfetch(w_password, (indexW+7)));
schedule8 = (uint32_t) (tex1Dfetch(w_password, (indexW+8)));
schedule9 = (uint32_t) (tex1Dfetch(w_password, (indexW+9)));
schedule10 = (uint32_t) (tex1Dfetch(w_password, (indexW+10)));
schedule11 = (uint32_t) (tex1Dfetch(w_password, (indexW+11)));
schedule12 = (uint32_t) (tex1Dfetch(w_password, (indexW+12)));
schedule13 = (uint32_t) (tex1Dfetch(w_password, (indexW+13)));
schedule14 = (uint32_t) (tex1Dfetch(w_password, (indexW+14)));
//Input password is shorter than FIRST_LENGHT
if(schedule14 == 0xFFFFFFFF) schedule14=0;
else if(method == MODE_USER_PASS) redo=1;
schedule15 = (uint32_t) (tex1Dfetch(w_password, (indexW+15)));
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;
//User password only
if(method == MODE_USER_PASS)
{
if(redo == 1)
{
schedule0 = (uint32_t) (tex1Dfetch(w_password, (indexW+16)));
schedule1 = (uint32_t) (tex1Dfetch(w_password, (indexW+17)));
schedule2 = (uint32_t) (tex1Dfetch(w_password, (indexW+18)));
schedule3 = (uint32_t) (tex1Dfetch(w_password, (indexW+19)));
schedule4 = (uint32_t) (tex1Dfetch(w_password, (indexW+20)));
schedule5 = (uint32_t) (tex1Dfetch(w_password, (indexW+21)));
schedule6 = (uint32_t) (tex1Dfetch(w_password, (indexW+22)));
schedule7 = (uint32_t) (tex1Dfetch(w_password, (indexW+23)));
schedule8 = (uint32_t) (tex1Dfetch(w_password, (indexW+24)));
schedule9 = (uint32_t) (tex1Dfetch(w_password, (indexW+25)));
schedule10 = (uint32_t) (tex1Dfetch(w_password, (indexW+26)));
schedule11 = (uint32_t) (tex1Dfetch(w_password, (indexW+27)));
schedule12 = (uint32_t) (tex1Dfetch(w_password, (indexW+28)));
schedule13 = (uint32_t) (tex1Dfetch(w_password, (indexW+29)));
schedule14 = (uint32_t) (tex1Dfetch(w_password, (indexW+30)));
schedule15 = (uint32_t) (tex1Dfetch(w_password, (indexW+31)));
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;
}
//----------------------------------------------------- SECOND HASH ------------------------------------------------
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;
}
//----------------------------------------------------- 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;
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);
schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6);
schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7);
hash0 ^= LOP3LUT_XOR(
LOP3LUT_XOR( (TS2[(hash7 >> 24) ] & 0x000000FF), (TS3[(hash7 >> 16) & 0xFF] & 0xFF000000), (TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000)),
(TS1[(hash7 ) & 0xFF] & 0x0000FF00), 0x01000000
); //RCON[0];
hash1 ^= hash0; hash2 ^= hash1; hash3 ^= hash2;
schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule4 >> 24], TS1[(schedule5 >> 16) & 0xFF], TS2[(schedule6 >> 8) & 0xFF]) , TS3[schedule7 & 0xFF] , hash0);
schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule5 >> 24], TS1[(schedule6 >> 16) & 0xFF], TS2[(schedule7 >> 8) & 0xFF]) , TS3[schedule4 & 0xFF] , hash1);
schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule6 >> 24], TS1[(schedule7 >> 16) & 0xFF], TS2[(schedule4 >> 8) & 0xFF]) , TS3[schedule5 & 0xFF] , hash2);
schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule7 >> 24], TS1[(schedule4 >> 16) & 0xFF], TS2[(schedule5 >> 8) & 0xFF]) , TS3[schedule6 & 0xFF] , hash3);
hash4 ^= (TS3[(hash3 >> 24) ] & 0xFF000000) ^
(TS0[(hash3 >> 16) & 0xFF] & 0x00FF0000) ^
(TS1[(hash3 >> 8) & 0xFF] & 0x0000FF00) ^
(TS2[(hash3 ) & 0xFF] & 0x000000FF);
hash5 ^= hash4;
hash6 ^= hash5;
hash7 ^= hash6;
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);
schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6);
schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7);
hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^
(TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^
(TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^
(TS1[(hash7 ) & 0xFF] & 0x0000FF00) ^ 0x02000000; //RCON[1];
hash1 ^= hash0; hash2 ^= hash1; hash3 ^= hash2;
schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule4 >> 24], TS1[(schedule5 >> 16) & 0xFF], TS2[(schedule6 >> 8) & 0xFF]) , TS3[schedule7 & 0xFF] , hash0);
schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule5 >> 24], TS1[(schedule6 >> 16) & 0xFF], TS2[(schedule7 >> 8) & 0xFF]) , TS3[schedule4 & 0xFF] , hash1);
schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule6 >> 24], TS1[(schedule7 >> 16) & 0xFF], TS2[(schedule4 >> 8) & 0xFF]) , TS3[schedule5 & 0xFF] , hash2);
schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule7 >> 24], TS1[(schedule4 >> 16) & 0xFF], TS2[(schedule5 >> 8) & 0xFF]) , TS3[schedule6 & 0xFF] , hash3);
hash4 ^= (TS3[(hash3 >> 24) ] & 0xFF000000) ^
(TS0[(hash3 >> 16) & 0xFF] & 0x00FF0000) ^
(TS1[(hash3 >> 8) & 0xFF] & 0x0000FF00) ^
(TS2[(hash3 ) & 0xFF] & 0x000000FF);
hash5 ^= hash4;
hash6 ^= hash5;
hash7 ^= hash6;
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);
schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6);
schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7);
hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^
(TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^
(TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^
(TS1[(hash7 ) & 0xFF] & 0x0000FF00) ^ 0x04000000; //RCON[2];
hash1 ^= hash0; hash2 ^= hash1; hash3 ^= hash2;
schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule4 >> 24], TS1[(schedule5 >> 16) & 0xFF], TS2[(schedule6 >> 8) & 0xFF]) , TS3[schedule7 & 0xFF] , hash0);
schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule5 >> 24], TS1[(schedule6 >> 16) & 0xFF], TS2[(schedule7 >> 8) & 0xFF]) , TS3[schedule4 & 0xFF] , hash1);
schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule6 >> 24], TS1[(schedule7 >> 16) & 0xFF], TS2[(schedule4 >> 8) & 0xFF]) , TS3[schedule5 & 0xFF] , hash2);
schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule7 >> 24], TS1[(schedule4 >> 16) & 0xFF], TS2[(schedule5 >> 8) & 0xFF]) , TS3[schedule6 & 0xFF] , hash3);
hash4 ^= (TS3[(hash3 >> 24) ] & 0xFF000000) ^
(TS0[(hash3 >> 16) & 0xFF] & 0x00FF0000) ^
(TS1[(hash3 >> 8) & 0xFF] & 0x0000FF00) ^
(TS2[(hash3 ) & 0xFF] & 0x000000FF);
hash5 ^= hash4;
hash6 ^= hash5;
hash7 ^= hash6;
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);
schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6);
schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7);
hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^
(TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^
(TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^
(TS1[(hash7 ) & 0xFF] & 0x0000FF00) ^ 0x08000000; //RCON[3];
hash1 ^= hash0; hash2 ^= hash1; hash3 ^= hash2;
schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule4 >> 24], TS1[(schedule5 >> 16) & 0xFF], TS2[(schedule6 >> 8) & 0xFF]) , TS3[schedule7 & 0xFF] , hash0);
schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule5 >> 24], TS1[(schedule6 >> 16) & 0xFF], TS2[(schedule7 >> 8) & 0xFF]) , TS3[schedule4 & 0xFF] , hash1);
schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule6 >> 24], TS1[(schedule7 >> 16) & 0xFF], TS2[(schedule4 >> 8) & 0xFF]) , TS3[schedule5 & 0xFF] , hash2);
schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule7 >> 24], TS1[(schedule4 >> 16) & 0xFF], TS2[(schedule5 >> 8) & 0xFF]) , TS3[schedule6 & 0xFF] , hash3);
hash4 ^= (TS3[(hash3 >> 24) ] & 0xFF000000) ^
(TS0[(hash3 >> 16) & 0xFF] & 0x00FF0000) ^
(TS1[(hash3 >> 8) & 0xFF] & 0x0000FF00) ^
(TS2[(hash3 ) & 0xFF] & 0x000000FF);
hash5 ^= hash4;
hash6 ^= hash5;
hash7 ^= hash6;
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);
schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6);
schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7);
hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^
(TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^
(TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^
(TS1[(hash7 ) & 0xFF] & 0x0000FF00) ^ 0x10000000; //RCON[4];
hash1 ^= hash0; hash2 ^= hash1; hash3 ^= hash2;
schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule4 >> 24], TS1[(schedule5 >> 16) & 0xFF], TS2[(schedule6 >> 8) & 0xFF]) , TS3[schedule7 & 0xFF] , hash0);
schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule5 >> 24], TS1[(schedule6 >> 16) & 0xFF], TS2[(schedule7 >> 8) & 0xFF]) , TS3[schedule4 & 0xFF] , hash1);
schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule6 >> 24], TS1[(schedule7 >> 16) & 0xFF], TS2[(schedule4 >> 8) & 0xFF]) , TS3[schedule5 & 0xFF] , hash2);
schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule7 >> 24], TS1[(schedule4 >> 16) & 0xFF], TS2[(schedule5 >> 8) & 0xFF]) , TS3[schedule6 & 0xFF] , hash3);
hash4 ^= (TS3[(hash3 >> 24) ] & 0xFF000000) ^
(TS0[(hash3 >> 16) & 0xFF] & 0x00FF0000) ^
(TS1[(hash3 >> 8) & 0xFF] & 0x0000FF00) ^
(TS2[(hash3 ) & 0xFF] & 0x000000FF);
hash5 ^= hash4;
hash6 ^= hash5;
hash7 ^= hash6;
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);
schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6);
schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7);
hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^
(TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^
(TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^
(TS1[(hash7 ) & 0xFF] & 0x0000FF00) ^ 0x20000000; //RCON[5];
hash1 ^= hash0; hash2 ^= hash1; hash3 ^= hash2;
schedule0 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule4 >> 24], TS1[(schedule5 >> 16) & 0xFF], TS2[(schedule6 >> 8) & 0xFF]) , TS3[schedule7 & 0xFF] , hash0);
schedule1 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule5 >> 24], TS1[(schedule6 >> 16) & 0xFF], TS2[(schedule7 >> 8) & 0xFF]) , TS3[schedule4 & 0xFF] , hash1);
schedule2 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule6 >> 24], TS1[(schedule7 >> 16) & 0xFF], TS2[(schedule4 >> 8) & 0xFF]) , TS3[schedule5 & 0xFF] , hash2);
schedule3 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule7 >> 24], TS1[(schedule4 >> 16) & 0xFF], TS2[(schedule5 >> 8) & 0xFF]) , TS3[schedule6 & 0xFF] , hash3);
hash4 ^= (TS3[(hash3 >> 24) ] & 0xFF000000) ^
(TS0[(hash3 >> 16) & 0xFF] & 0x00FF0000) ^
(TS1[(hash3 >> 8) & 0xFF] & 0x0000FF00) ^
(TS2[(hash3 ) & 0xFF] & 0x000000FF);
hash5 ^= hash4;
hash6 ^= hash5;
hash7 ^= hash6;
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);
schedule6 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule2 >> 24], TS1[(schedule3 >> 16) & 0xFF], TS2[(schedule0 >> 8) & 0xFF]) , TS3[schedule1 & 0xFF] , hash6);
schedule7 = LOP3LUT_XOR(LOP3LUT_XOR(TS0[schedule3 >> 24], TS1[(schedule0 >> 16) & 0xFF], TS2[(schedule1 >> 8) & 0xFF]) , TS3[schedule2 & 0xFF] , hash7);
hash0 ^= (TS2[(hash7 >> 24) ] & 0x000000FF) ^
(TS3[(hash7 >> 16) & 0xFF] & 0xFF000000) ^
(TS0[(hash7 >> 8) & 0xFF] & 0x00FF0000) ^
(TS1[(hash7 ) & 0xFF] & 0x0000FF00) ^ 0x40000000; //RCON[6];
hash1 ^= hash0;
hash2 ^= hash1;
hash3 ^= hash2;
schedule0 = (TS2[(schedule4 >> 24) ] & 0xFF000000) ^
(TS3[(schedule5 >> 16) & 0xFF] & 0x00FF0000) ^
(TS0[(schedule6 >> 8) & 0xFF] & 0x0000FF00) ^
(TS1[(schedule7 ) & 0xFF] & 0x000000FF) ^ hash0;
schedule1 = (TS2[(schedule5 >> 24) ] & 0xFF000000) ^
(TS3[(schedule6 >> 16) & 0xFF] & 0x00FF0000) ^
(TS0[(schedule7 >> 8) & 0xFF] & 0x0000FF00) ^
(TS1[(schedule4 ) & 0xFF] & 0x000000FF) ^ hash1;
schedule2 = (TS2[(schedule6 >> 24) ] & 0xFF000000) ^
(TS3[(schedule7 >> 16) & 0xFF] & 0x00FF0000) ^
(TS0[(schedule4 >> 8) & 0xFF] & 0x0000FF00) ^
(TS1[(schedule5 ) & 0xFF] & 0x000000FF) ^ hash2;
schedule3 = (TS2[(schedule7 >> 24) ] & 0xFF000000) ^
(TS3[(schedule4 >> 16) & 0xFF] & 0x00FF0000) ^
(TS0[(schedule5 >> 8) & 0xFF] & 0x0000FF00) ^
(TS1[(schedule6 ) & 0xFF] & 0x000000FF) ^ hash3;
schedule4 = __byte_perm(schedule0, 0, 0x0123);
schedule5 = __byte_perm(schedule1, 0, 0x0123);
schedule6 = __byte_perm(schedule2, 0, 0x0123);
schedule7 = __byte_perm(schedule3, 0, 0x0123);
if (
((vmkKey[0] ^ ((uint8_t) schedule4)) == 0x2c) &&
((vmkKey[1] ^ ((uint8_t) (schedule4 >> 8))) == 0x00) &&
((vmkKey[4] ^ ((uint8_t) schedule5)) == 0x01) &&
((vmkKey[5] ^ ((uint8_t) (schedule5 >> 8))) == 0x00) &&
((vmkKey[9] ^ ((uint8_t) (schedule6 >> 8))) == 0x20)
)
{
if(
(strict_check == 0 && ((vmkKey[8] ^ ((uint8_t) schedule6)) <= 0x05))
||
(strict_check == 1 && ((vmkKey[8] ^ ((uint8_t) schedule6)) == 0x03))
)
{
*found = gIndex;
break;
}
}
gIndex += (blockDim.x * gridDim.x);
}
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 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,
int method
)
{
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*PSW_INT_SIZE);
int8_t redo=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*PSW_INT_SIZE);
redo=0;
schedule0 = (uint32_t) (tex1Dfetch(w_password, (indexW+0)));
schedule1 = (uint32_t) (tex1Dfetch(w_password, (indexW+1)));
schedule2 = (uint32_t) (tex1Dfetch(w_password, (indexW+2)));
schedule3 = (uint32_t) (tex1Dfetch(w_password, (indexW+3)));
schedule4 = (uint32_t) (tex1Dfetch(w_password, (indexW+4)));
schedule5 = (uint32_t) (tex1Dfetch(w_password, (indexW+5)));
schedule6 = (uint32_t) (tex1Dfetch(w_password, (indexW+6)));
schedule7 = (uint32_t) (tex1Dfetch(w_password, (indexW+7)));
schedule8 = (uint32_t) (tex1Dfetch(w_password, (indexW+8)));
schedule9 = (uint32_t) (tex1Dfetch(w_password, (indexW+9)));
schedule10 = (uint32_t) (tex1Dfetch(w_password, (indexW+10)));
schedule11 = (uint32_t) (tex1Dfetch(w_password, (indexW+11)));
schedule12 = (uint32_t) (tex1Dfetch(w_password, (indexW+12)));
schedule13 = (uint32_t) (tex1Dfetch(w_password, (indexW+13)));
schedule14 = (uint32_t) (tex1Dfetch(w_password, (indexW+14)));
//Input password is shorter than FIRST_LENGHT
if(schedule14 == 0xFFFFFFFF) schedule14=0;
else if(method == MODE_USER_PASS) redo=1;
schedule15 = (uint32_t) (tex1Dfetch(w_password, (indexW+15)));
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;
//User password only
if(method == MODE_USER_PASS)
{
if(redo == 1)
{
schedule0 = (uint32_t) (tex1Dfetch(w_password, (indexW+16)));
schedule1 = (uint32_t) (tex1Dfetch(w_password, (indexW+17)));
schedule2 = (uint32_t) (tex1Dfetch(w_password, (indexW+18)));
schedule3 = (uint32_t) (tex1Dfetch(w_password, (indexW+19)));
schedule4 = (uint32_t) (tex1Dfetch(w_password, (indexW+20)));
schedule5 = (uint32_t) (tex1Dfetch(w_password, (indexW+21)));
schedule6 = (uint32_t) (tex1Dfetch(w_password, (indexW+22)));
schedule7 = (uint32_t) (tex1Dfetch(w_password, (indexW+23)));
schedule8 = (uint32_t) (tex1Dfetch(w_password, (indexW+24)));
schedule9 = (uint32_t) (tex1Dfetch(w_password, (indexW+25)));
schedule10 = (uint32_t) (tex1Dfetch(w_password, (indexW+26)));
schedule11 = (uint32_t) (tex1Dfetch(w_password, (indexW+27)));
schedule12 = (uint32_t) (tex1Dfetch(w_password, (indexW+28)));
schedule13 = (uint32_t) (tex1Dfetch(w_password, (indexW+29)));
schedule14 = (uint32_t) (tex1Dfetch(w_password, (indexW+30)));
schedule15 = (uint32_t) (tex1Dfetch(w_password, (indexW+31)));
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;
}
//----------------------------------------------------- SECOND HASH ------------------------------------------------
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;
}
Reference in New Issue View Git Blame Copy Permalink