From d537712f2715e4d8e772f8a1cdd351a73e6a2c20 Mon Sep 17 00:00:00 2001 From: Fist0urs Date: Sat, 6 May 2017 16:55:36 +0200 Subject: [PATCH] Both DPAPImk v1 and v2 work for single hash, still a bug on multi-hash Remaining this bug + tests.pl before PR --- OpenCL/inc_types.cl | 13 +- OpenCL/m15300.cl | 1371 ++++++++++++++++++++++++++++++++++++------- include/interface.h | 24 +- src/interface.c | 141 ++--- 4 files changed, 1216 insertions(+), 333 deletions(-) diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl index 7360c3019..fae6d9165 100644 --- a/OpenCL/inc_types.cl +++ b/OpenCL/inc_types.cl @@ -1022,14 +1022,16 @@ typedef struct } keepass_t; -/* Fist0urs */ typedef struct { u32 version; u32 context; u8 SID_tmp[64]; - u32 SID[64]; + u32 SID[32]; + u32 SID_len; + u32 SID_offset; + u32 userKey[5]; char cipher_algo[16]; char hash_algo[16]; @@ -1038,7 +1040,6 @@ typedef struct u32 contents[128]; } dpapimk_t; -/* Fist0urs_end */ typedef struct { @@ -1379,14 +1380,13 @@ typedef struct } keepass_tmp_t; -/* Fist0urs */ typedef struct { /* dedicated to hmac-sha1 */ u32 ipad[5]; u32 opad[5]; - u32 dgst[5]; - u32 out[5]; + u32 dgst[10]; + u32 out[10]; /* dedicated to hmac-sha512 */ u64 ipad64[8]; @@ -1395,7 +1395,6 @@ typedef struct u64 out64[16]; } dpapimk_tmp_t; -/* Fist0urs_end */ typedef struct { diff --git a/OpenCL/m15300.cl b/OpenCL/m15300.cl index 05a46e74b..e0fb11dc0 100644 --- a/OpenCL/m15300.cl +++ b/OpenCL/m15300.cl @@ -11,72 +11,555 @@ #include "inc_types.cl" #include "inc_common.cl" #include "inc_simd.cl" - +#include "inc_rp.h" +#include "inc_rp.cl" #include "inc_cipher_aes.cl" -/* Fist0urs */ -void u32_to_hex_lower (const u32 v, u8 hex[8]) -{ - const u8 tbl[0x10] = - { - '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', - 'a', 'b', 'c', 'd', 'e', 'f', - }; - - hex[1] = tbl[v >> 0 & 15]; - hex[0] = tbl[v >> 4 & 15]; - hex[3] = tbl[v >> 8 & 15]; - hex[2] = tbl[v >> 12 & 15]; - hex[5] = tbl[v >> 16 & 15]; - hex[4] = tbl[v >> 20 & 15]; - hex[7] = tbl[v >> 24 & 15]; - hex[6] = tbl[v >> 28 & 15]; +#define PERM_OP(a,b,tt,n,m) \ +{ \ + tt = a >> n; \ + tt = tt ^ b; \ + tt = tt & m; \ + b = b ^ tt; \ + tt = tt << n; \ + a = a ^ tt; \ } -void u64_to_hex_lower (const u64 v, u8 hex[16]) -{ - const u8 tbl[0x10] = - { - '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', - 'a', 'b', 'c', 'd', 'e', 'f', - }; - - hex[ 1] = tbl[v >> 0 & 15]; - hex[ 0] = tbl[v >> 4 & 15]; - hex[ 3] = tbl[v >> 8 & 15]; - hex[ 2] = tbl[v >> 12 & 15]; - hex[ 5] = tbl[v >> 16 & 15]; - hex[ 4] = tbl[v >> 20 & 15]; - hex[ 7] = tbl[v >> 24 & 15]; - hex[ 6] = tbl[v >> 28 & 15]; - hex[ 9] = tbl[v >> 32 & 15]; - hex[ 8] = tbl[v >> 36 & 15]; - hex[11] = tbl[v >> 40 & 15]; - hex[10] = tbl[v >> 44 & 15]; - hex[13] = tbl[v >> 48 & 15]; - hex[12] = tbl[v >> 52 & 15]; - hex[15] = tbl[v >> 56 & 15]; - hex[14] = tbl[v >> 60 & 15]; +#define HPERM_OP(a,tt,n,m) \ +{ \ + tt = a << (16 + n); \ + tt = tt ^ a; \ + tt = tt & m; \ + a = a ^ tt; \ + tt = tt >> (16 + n); \ + a = a ^ tt; \ } -int -pretty_print(char *message, void *data, int len) +#define IP(l,r,tt) \ +{ \ + PERM_OP (r, l, tt, 4, 0x0f0f0f0f); \ + PERM_OP (l, r, tt, 16, 0x0000ffff); \ + PERM_OP (r, l, tt, 2, 0x33333333); \ + PERM_OP (l, r, tt, 8, 0x00ff00ff); \ + PERM_OP (r, l, tt, 1, 0x55555555); \ +} + +#define FP(l,r,tt) \ +{ \ + PERM_OP (l, r, tt, 1, 0x55555555); \ + PERM_OP (r, l, tt, 8, 0x00ff00ff); \ + PERM_OP (l, r, tt, 2, 0x33333333); \ + PERM_OP (r, l, tt, 16, 0x0000ffff); \ + PERM_OP (l, r, tt, 4, 0x0f0f0f0f); \ +} + +__constant u32a c_SPtrans[8][64] = { - int g = 0; - for (int i = 0 ; i < len; i++) { - if (g == 0) - { - printf("%s: ", message); - g++; - } - printf("%02x", ((char *)(data))[i]); + /* nibble 0 */ + 0x02080800, 0x00080000, 0x02000002, 0x02080802, + 0x02000000, 0x00080802, 0x00080002, 0x02000002, + 0x00080802, 0x02080800, 0x02080000, 0x00000802, + 0x02000802, 0x02000000, 0x00000000, 0x00080002, + 0x00080000, 0x00000002, 0x02000800, 0x00080800, + 0x02080802, 0x02080000, 0x00000802, 0x02000800, + 0x00000002, 0x00000800, 0x00080800, 0x02080002, + 0x00000800, 0x02000802, 0x02080002, 0x00000000, + 0x00000000, 0x02080802, 0x02000800, 0x00080002, + 0x02080800, 0x00080000, 0x00000802, 0x02000800, + 0x02080002, 0x00000800, 0x00080800, 0x02000002, + 0x00080802, 0x00000002, 0x02000002, 0x02080000, + 0x02080802, 0x00080800, 0x02080000, 0x02000802, + 0x02000000, 0x00000802, 0x00080002, 0x00000000, + 0x00080000, 0x02000000, 0x02000802, 0x02080800, + 0x00000002, 0x02080002, 0x00000800, 0x00080802, + }, + { + /* nibble 1 */ + 0x40108010, 0x00000000, 0x00108000, 0x40100000, + 0x40000010, 0x00008010, 0x40008000, 0x00108000, + 0x00008000, 0x40100010, 0x00000010, 0x40008000, + 0x00100010, 0x40108000, 0x40100000, 0x00000010, + 0x00100000, 0x40008010, 0x40100010, 0x00008000, + 0x00108010, 0x40000000, 0x00000000, 0x00100010, + 0x40008010, 0x00108010, 0x40108000, 0x40000010, + 0x40000000, 0x00100000, 0x00008010, 0x40108010, + 0x00100010, 0x40108000, 0x40008000, 0x00108010, + 0x40108010, 0x00100010, 0x40000010, 0x00000000, + 0x40000000, 0x00008010, 0x00100000, 0x40100010, + 0x00008000, 0x40000000, 0x00108010, 0x40008010, + 0x40108000, 0x00008000, 0x00000000, 0x40000010, + 0x00000010, 0x40108010, 0x00108000, 0x40100000, + 0x40100010, 0x00100000, 0x00008010, 0x40008000, + 0x40008010, 0x00000010, 0x40100000, 0x00108000, + }, + { + /* nibble 2 */ + 0x04000001, 0x04040100, 0x00000100, 0x04000101, + 0x00040001, 0x04000000, 0x04000101, 0x00040100, + 0x04000100, 0x00040000, 0x04040000, 0x00000001, + 0x04040101, 0x00000101, 0x00000001, 0x04040001, + 0x00000000, 0x00040001, 0x04040100, 0x00000100, + 0x00000101, 0x04040101, 0x00040000, 0x04000001, + 0x04040001, 0x04000100, 0x00040101, 0x04040000, + 0x00040100, 0x00000000, 0x04000000, 0x00040101, + 0x04040100, 0x00000100, 0x00000001, 0x00040000, + 0x00000101, 0x00040001, 0x04040000, 0x04000101, + 0x00000000, 0x04040100, 0x00040100, 0x04040001, + 0x00040001, 0x04000000, 0x04040101, 0x00000001, + 0x00040101, 0x04000001, 0x04000000, 0x04040101, + 0x00040000, 0x04000100, 0x04000101, 0x00040100, + 0x04000100, 0x00000000, 0x04040001, 0x00000101, + 0x04000001, 0x00040101, 0x00000100, 0x04040000, + }, + { + /* nibble 3 */ + 0x00401008, 0x10001000, 0x00000008, 0x10401008, + 0x00000000, 0x10400000, 0x10001008, 0x00400008, + 0x10401000, 0x10000008, 0x10000000, 0x00001008, + 0x10000008, 0x00401008, 0x00400000, 0x10000000, + 0x10400008, 0x00401000, 0x00001000, 0x00000008, + 0x00401000, 0x10001008, 0x10400000, 0x00001000, + 0x00001008, 0x00000000, 0x00400008, 0x10401000, + 0x10001000, 0x10400008, 0x10401008, 0x00400000, + 0x10400008, 0x00001008, 0x00400000, 0x10000008, + 0x00401000, 0x10001000, 0x00000008, 0x10400000, + 0x10001008, 0x00000000, 0x00001000, 0x00400008, + 0x00000000, 0x10400008, 0x10401000, 0x00001000, + 0x10000000, 0x10401008, 0x00401008, 0x00400000, + 0x10401008, 0x00000008, 0x10001000, 0x00401008, + 0x00400008, 0x00401000, 0x10400000, 0x10001008, + 0x00001008, 0x10000000, 0x10000008, 0x10401000, + }, + { + /* nibble 4 */ + 0x08000000, 0x00010000, 0x00000400, 0x08010420, + 0x08010020, 0x08000400, 0x00010420, 0x08010000, + 0x00010000, 0x00000020, 0x08000020, 0x00010400, + 0x08000420, 0x08010020, 0x08010400, 0x00000000, + 0x00010400, 0x08000000, 0x00010020, 0x00000420, + 0x08000400, 0x00010420, 0x00000000, 0x08000020, + 0x00000020, 0x08000420, 0x08010420, 0x00010020, + 0x08010000, 0x00000400, 0x00000420, 0x08010400, + 0x08010400, 0x08000420, 0x00010020, 0x08010000, + 0x00010000, 0x00000020, 0x08000020, 0x08000400, + 0x08000000, 0x00010400, 0x08010420, 0x00000000, + 0x00010420, 0x08000000, 0x00000400, 0x00010020, + 0x08000420, 0x00000400, 0x00000000, 0x08010420, + 0x08010020, 0x08010400, 0x00000420, 0x00010000, + 0x00010400, 0x08010020, 0x08000400, 0x00000420, + 0x00000020, 0x00010420, 0x08010000, 0x08000020, + }, + { + /* nibble 5 */ + 0x80000040, 0x00200040, 0x00000000, 0x80202000, + 0x00200040, 0x00002000, 0x80002040, 0x00200000, + 0x00002040, 0x80202040, 0x00202000, 0x80000000, + 0x80002000, 0x80000040, 0x80200000, 0x00202040, + 0x00200000, 0x80002040, 0x80200040, 0x00000000, + 0x00002000, 0x00000040, 0x80202000, 0x80200040, + 0x80202040, 0x80200000, 0x80000000, 0x00002040, + 0x00000040, 0x00202000, 0x00202040, 0x80002000, + 0x00002040, 0x80000000, 0x80002000, 0x00202040, + 0x80202000, 0x00200040, 0x00000000, 0x80002000, + 0x80000000, 0x00002000, 0x80200040, 0x00200000, + 0x00200040, 0x80202040, 0x00202000, 0x00000040, + 0x80202040, 0x00202000, 0x00200000, 0x80002040, + 0x80000040, 0x80200000, 0x00202040, 0x00000000, + 0x00002000, 0x80000040, 0x80002040, 0x80202000, + 0x80200000, 0x00002040, 0x00000040, 0x80200040, + }, + { + /* nibble 6 */ + 0x00004000, 0x00000200, 0x01000200, 0x01000004, + 0x01004204, 0x00004004, 0x00004200, 0x00000000, + 0x01000000, 0x01000204, 0x00000204, 0x01004000, + 0x00000004, 0x01004200, 0x01004000, 0x00000204, + 0x01000204, 0x00004000, 0x00004004, 0x01004204, + 0x00000000, 0x01000200, 0x01000004, 0x00004200, + 0x01004004, 0x00004204, 0x01004200, 0x00000004, + 0x00004204, 0x01004004, 0x00000200, 0x01000000, + 0x00004204, 0x01004000, 0x01004004, 0x00000204, + 0x00004000, 0x00000200, 0x01000000, 0x01004004, + 0x01000204, 0x00004204, 0x00004200, 0x00000000, + 0x00000200, 0x01000004, 0x00000004, 0x01000200, + 0x00000000, 0x01000204, 0x01000200, 0x00004200, + 0x00000204, 0x00004000, 0x01004204, 0x01000000, + 0x01004200, 0x00000004, 0x00004004, 0x01004204, + 0x01000004, 0x01004200, 0x01004000, 0x00004004, + }, + { + /* nibble 7 */ + 0x20800080, 0x20820000, 0x00020080, 0x00000000, + 0x20020000, 0x00800080, 0x20800000, 0x20820080, + 0x00000080, 0x20000000, 0x00820000, 0x00020080, + 0x00820080, 0x20020080, 0x20000080, 0x20800000, + 0x00020000, 0x00820080, 0x00800080, 0x20020000, + 0x20820080, 0x20000080, 0x00000000, 0x00820000, + 0x20000000, 0x00800000, 0x20020080, 0x20800080, + 0x00800000, 0x00020000, 0x20820000, 0x00000080, + 0x00800000, 0x00020000, 0x20000080, 0x20820080, + 0x00020080, 0x20000000, 0x00000000, 0x00820000, + 0x20800080, 0x20020080, 0x20020000, 0x00800080, + 0x20820000, 0x00000080, 0x00800080, 0x20020000, + 0x20820080, 0x00800000, 0x20800000, 0x20000080, + 0x00820000, 0x00020080, 0x20020080, 0x20800000, + 0x00000080, 0x20820000, 0x00820080, 0x00000000, + 0x20000000, 0x20800080, 0x00020000, 0x00820080, + }, +}; + +__constant u32a c_skb[8][64] = +{ + { + 0x00000000, 0x00000010, 0x20000000, 0x20000010, + 0x00010000, 0x00010010, 0x20010000, 0x20010010, + 0x00000800, 0x00000810, 0x20000800, 0x20000810, + 0x00010800, 0x00010810, 0x20010800, 0x20010810, + 0x00000020, 0x00000030, 0x20000020, 0x20000030, + 0x00010020, 0x00010030, 0x20010020, 0x20010030, + 0x00000820, 0x00000830, 0x20000820, 0x20000830, + 0x00010820, 0x00010830, 0x20010820, 0x20010830, + 0x00080000, 0x00080010, 0x20080000, 0x20080010, + 0x00090000, 0x00090010, 0x20090000, 0x20090010, + 0x00080800, 0x00080810, 0x20080800, 0x20080810, + 0x00090800, 0x00090810, 0x20090800, 0x20090810, + 0x00080020, 0x00080030, 0x20080020, 0x20080030, + 0x00090020, 0x00090030, 0x20090020, 0x20090030, + 0x00080820, 0x00080830, 0x20080820, 0x20080830, + 0x00090820, 0x00090830, 0x20090820, 0x20090830, + }, + { + 0x00000000, 0x02000000, 0x00002000, 0x02002000, + 0x00200000, 0x02200000, 0x00202000, 0x02202000, + 0x00000004, 0x02000004, 0x00002004, 0x02002004, + 0x00200004, 0x02200004, 0x00202004, 0x02202004, + 0x00000400, 0x02000400, 0x00002400, 0x02002400, + 0x00200400, 0x02200400, 0x00202400, 0x02202400, + 0x00000404, 0x02000404, 0x00002404, 0x02002404, + 0x00200404, 0x02200404, 0x00202404, 0x02202404, + 0x10000000, 0x12000000, 0x10002000, 0x12002000, + 0x10200000, 0x12200000, 0x10202000, 0x12202000, + 0x10000004, 0x12000004, 0x10002004, 0x12002004, + 0x10200004, 0x12200004, 0x10202004, 0x12202004, + 0x10000400, 0x12000400, 0x10002400, 0x12002400, + 0x10200400, 0x12200400, 0x10202400, 0x12202400, + 0x10000404, 0x12000404, 0x10002404, 0x12002404, + 0x10200404, 0x12200404, 0x10202404, 0x12202404, + }, + { + 0x00000000, 0x00000001, 0x00040000, 0x00040001, + 0x01000000, 0x01000001, 0x01040000, 0x01040001, + 0x00000002, 0x00000003, 0x00040002, 0x00040003, + 0x01000002, 0x01000003, 0x01040002, 0x01040003, + 0x00000200, 0x00000201, 0x00040200, 0x00040201, + 0x01000200, 0x01000201, 0x01040200, 0x01040201, + 0x00000202, 0x00000203, 0x00040202, 0x00040203, + 0x01000202, 0x01000203, 0x01040202, 0x01040203, + 0x08000000, 0x08000001, 0x08040000, 0x08040001, + 0x09000000, 0x09000001, 0x09040000, 0x09040001, + 0x08000002, 0x08000003, 0x08040002, 0x08040003, + 0x09000002, 0x09000003, 0x09040002, 0x09040003, + 0x08000200, 0x08000201, 0x08040200, 0x08040201, + 0x09000200, 0x09000201, 0x09040200, 0x09040201, + 0x08000202, 0x08000203, 0x08040202, 0x08040203, + 0x09000202, 0x09000203, 0x09040202, 0x09040203, + }, + { + 0x00000000, 0x00100000, 0x00000100, 0x00100100, + 0x00000008, 0x00100008, 0x00000108, 0x00100108, + 0x00001000, 0x00101000, 0x00001100, 0x00101100, + 0x00001008, 0x00101008, 0x00001108, 0x00101108, + 0x04000000, 0x04100000, 0x04000100, 0x04100100, + 0x04000008, 0x04100008, 0x04000108, 0x04100108, + 0x04001000, 0x04101000, 0x04001100, 0x04101100, + 0x04001008, 0x04101008, 0x04001108, 0x04101108, + 0x00020000, 0x00120000, 0x00020100, 0x00120100, + 0x00020008, 0x00120008, 0x00020108, 0x00120108, + 0x00021000, 0x00121000, 0x00021100, 0x00121100, + 0x00021008, 0x00121008, 0x00021108, 0x00121108, + 0x04020000, 0x04120000, 0x04020100, 0x04120100, + 0x04020008, 0x04120008, 0x04020108, 0x04120108, + 0x04021000, 0x04121000, 0x04021100, 0x04121100, + 0x04021008, 0x04121008, 0x04021108, 0x04121108, + }, + { + 0x00000000, 0x10000000, 0x00010000, 0x10010000, + 0x00000004, 0x10000004, 0x00010004, 0x10010004, + 0x20000000, 0x30000000, 0x20010000, 0x30010000, + 0x20000004, 0x30000004, 0x20010004, 0x30010004, + 0x00100000, 0x10100000, 0x00110000, 0x10110000, + 0x00100004, 0x10100004, 0x00110004, 0x10110004, + 0x20100000, 0x30100000, 0x20110000, 0x30110000, + 0x20100004, 0x30100004, 0x20110004, 0x30110004, + 0x00001000, 0x10001000, 0x00011000, 0x10011000, + 0x00001004, 0x10001004, 0x00011004, 0x10011004, + 0x20001000, 0x30001000, 0x20011000, 0x30011000, + 0x20001004, 0x30001004, 0x20011004, 0x30011004, + 0x00101000, 0x10101000, 0x00111000, 0x10111000, + 0x00101004, 0x10101004, 0x00111004, 0x10111004, + 0x20101000, 0x30101000, 0x20111000, 0x30111000, + 0x20101004, 0x30101004, 0x20111004, 0x30111004, + }, + { + 0x00000000, 0x08000000, 0x00000008, 0x08000008, + 0x00000400, 0x08000400, 0x00000408, 0x08000408, + 0x00020000, 0x08020000, 0x00020008, 0x08020008, + 0x00020400, 0x08020400, 0x00020408, 0x08020408, + 0x00000001, 0x08000001, 0x00000009, 0x08000009, + 0x00000401, 0x08000401, 0x00000409, 0x08000409, + 0x00020001, 0x08020001, 0x00020009, 0x08020009, + 0x00020401, 0x08020401, 0x00020409, 0x08020409, + 0x02000000, 0x0A000000, 0x02000008, 0x0A000008, + 0x02000400, 0x0A000400, 0x02000408, 0x0A000408, + 0x02020000, 0x0A020000, 0x02020008, 0x0A020008, + 0x02020400, 0x0A020400, 0x02020408, 0x0A020408, + 0x02000001, 0x0A000001, 0x02000009, 0x0A000009, + 0x02000401, 0x0A000401, 0x02000409, 0x0A000409, + 0x02020001, 0x0A020001, 0x02020009, 0x0A020009, + 0x02020401, 0x0A020401, 0x02020409, 0x0A020409, + }, + { + 0x00000000, 0x00000100, 0x00080000, 0x00080100, + 0x01000000, 0x01000100, 0x01080000, 0x01080100, + 0x00000010, 0x00000110, 0x00080010, 0x00080110, + 0x01000010, 0x01000110, 0x01080010, 0x01080110, + 0x00200000, 0x00200100, 0x00280000, 0x00280100, + 0x01200000, 0x01200100, 0x01280000, 0x01280100, + 0x00200010, 0x00200110, 0x00280010, 0x00280110, + 0x01200010, 0x01200110, 0x01280010, 0x01280110, + 0x00000200, 0x00000300, 0x00080200, 0x00080300, + 0x01000200, 0x01000300, 0x01080200, 0x01080300, + 0x00000210, 0x00000310, 0x00080210, 0x00080310, + 0x01000210, 0x01000310, 0x01080210, 0x01080310, + 0x00200200, 0x00200300, 0x00280200, 0x00280300, + 0x01200200, 0x01200300, 0x01280200, 0x01280300, + 0x00200210, 0x00200310, 0x00280210, 0x00280310, + 0x01200210, 0x01200310, 0x01280210, 0x01280310, + }, + { + 0x00000000, 0x04000000, 0x00040000, 0x04040000, + 0x00000002, 0x04000002, 0x00040002, 0x04040002, + 0x00002000, 0x04002000, 0x00042000, 0x04042000, + 0x00002002, 0x04002002, 0x00042002, 0x04042002, + 0x00000020, 0x04000020, 0x00040020, 0x04040020, + 0x00000022, 0x04000022, 0x00040022, 0x04040022, + 0x00002020, 0x04002020, 0x00042020, 0x04042020, + 0x00002022, 0x04002022, 0x00042022, 0x04042022, + 0x00000800, 0x04000800, 0x00040800, 0x04040800, + 0x00000802, 0x04000802, 0x00040802, 0x04040802, + 0x00002800, 0x04002800, 0x00042800, 0x04042800, + 0x00002802, 0x04002802, 0x00042802, 0x04042802, + 0x00000820, 0x04000820, 0x00040820, 0x04040820, + 0x00000822, 0x04000822, 0x00040822, 0x04040822, + 0x00002820, 0x04002820, 0x00042820, 0x04042820, + 0x00002822, 0x04002822, 0x00042822, 0x04042822 } - printf("\n"); +}; - return 1; +#if VECT_SIZE == 1 +#define BOX(i,n,S) (S)[(n)][(i)] +#elif VECT_SIZE == 2 +#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1]) +#elif VECT_SIZE == 4 +#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3]) +#elif VECT_SIZE == 8 +#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7]) +#elif VECT_SIZE == 16 +#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf]) +#endif + +#if VECT_SIZE == 1 +#define BOX1(i,S) (S)[(i)] +#elif VECT_SIZE == 2 +#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1]) +#elif VECT_SIZE == 4 +#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3]) +#elif VECT_SIZE == 8 +#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7]) +#elif VECT_SIZE == 16 +#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7], (S)[(i).s8], (S)[(i).s9], (S)[(i).sa], (S)[(i).sb], (S)[(i).sc], (S)[(i).sd], (S)[(i).se], (S)[(i).sf]) +#endif + +void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64]) +{ + u32x tt; + + u32x r = data[0]; + u32x l = data[1]; + + IP (r, l, tt); + + r = rotl32 (r, 3u); + l = rotl32 (l, 3u); + + #ifdef _unroll + #pragma unroll + #endif + for (u32 i = 0; i < 16; i += 2) + { + u32x u; + u32x t; + + u = Kc[i + 0] ^ r; + t = Kd[i + 0] ^ rotl32 (r, 28u); + + l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) + | BOX (((u >> 10) & 0x3f), 2, s_SPtrans) + | BOX (((u >> 18) & 0x3f), 4, s_SPtrans) + | BOX (((u >> 26) & 0x3f), 6, s_SPtrans) + | BOX (((t >> 2) & 0x3f), 1, s_SPtrans) + | BOX (((t >> 10) & 0x3f), 3, s_SPtrans) + | BOX (((t >> 18) & 0x3f), 5, s_SPtrans) + | BOX (((t >> 26) & 0x3f), 7, s_SPtrans); + + u = Kc[i + 1] ^ l; + t = Kd[i + 1] ^ rotl32 (l, 28u); + + r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) + | BOX (((u >> 10) & 0x3f), 2, s_SPtrans) + | BOX (((u >> 18) & 0x3f), 4, s_SPtrans) + | BOX (((u >> 26) & 0x3f), 6, s_SPtrans) + | BOX (((t >> 2) & 0x3f), 1, s_SPtrans) + | BOX (((t >> 10) & 0x3f), 3, s_SPtrans) + | BOX (((t >> 18) & 0x3f), 5, s_SPtrans) + | BOX (((t >> 26) & 0x3f), 7, s_SPtrans); + } + + l = rotl32 (l, 29u); + r = rotl32 (r, 29u); + + FP (r, l, tt); + + iv[0] = l; + iv[1] = r; +} + +void _des_crypt_decrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64]) +{ + u32x tt; + + u32x r = data[0]; + u32x l = data[1]; + + IP (r, l, tt); + + r = rotl32 (r, 3u); + l = rotl32 (l, 3u); + + #ifdef _unroll + #pragma unroll + #endif + for (u32 i = 16; i > 0; i -= 2) + { + u32x u; + u32x t; + + u = Kc[i - 1] ^ r; + t = Kd[i - 1] ^ rotl32 (r, 28u); + + l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) + | BOX (((u >> 10) & 0x3f), 2, s_SPtrans) + | BOX (((u >> 18) & 0x3f), 4, s_SPtrans) + | BOX (((u >> 26) & 0x3f), 6, s_SPtrans) + | BOX (((t >> 2) & 0x3f), 1, s_SPtrans) + | BOX (((t >> 10) & 0x3f), 3, s_SPtrans) + | BOX (((t >> 18) & 0x3f), 5, s_SPtrans) + | BOX (((t >> 26) & 0x3f), 7, s_SPtrans); + + u = Kc[i - 2] ^ l; + t = Kd[i - 2] ^ rotl32 (l, 28u); + + r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans) + | BOX (((u >> 10) & 0x3f), 2, s_SPtrans) + | BOX (((u >> 18) & 0x3f), 4, s_SPtrans) + | BOX (((u >> 26) & 0x3f), 6, s_SPtrans) + | BOX (((t >> 2) & 0x3f), 1, s_SPtrans) + | BOX (((t >> 10) & 0x3f), 3, s_SPtrans) + | BOX (((t >> 18) & 0x3f), 5, s_SPtrans) + | BOX (((t >> 26) & 0x3f), 7, s_SPtrans); + } + + l = rotl32 (l, 29u); + r = rotl32 (r, 29u); + + FP (r, l, tt); + + iv[0] = l; + iv[1] = r; +} + +void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64]) +{ + u32x tt; + + PERM_OP (d, c, tt, 4, 0x0f0f0f0f); + HPERM_OP (c, tt, 2, 0xcccc0000); + HPERM_OP (d, tt, 2, 0xcccc0000); + PERM_OP (d, c, tt, 1, 0x55555555); + PERM_OP (c, d, tt, 8, 0x00ff00ff); + PERM_OP (d, c, tt, 1, 0x55555555); + + d = ((d & 0x000000ff) << 16) + | ((d & 0x0000ff00) << 0) + | ((d & 0x00ff0000) >> 16) + | ((c & 0xf0000000) >> 4); + + c = c & 0x0fffffff; + + #ifdef _unroll + #pragma unroll + #endif + for (u32 i = 0; i < 16; i++) + { + if ((i < 2) || (i == 8) || (i == 15)) + { + c = ((c >> 1) | (c << 27)); + d = ((d >> 1) | (d << 27)); + } + else + { + c = ((c >> 2) | (c << 26)); + d = ((d >> 2) | (d << 26)); + } + + c = c & 0x0fffffff; + d = d & 0x0fffffff; + + const u32x c00 = (c >> 0) & 0x0000003f; + const u32x c06 = (c >> 6) & 0x00383003; + const u32x c07 = (c >> 7) & 0x0000003c; + const u32x c13 = (c >> 13) & 0x0000060f; + const u32x c20 = (c >> 20) & 0x00000001; + + u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb) + | BOX (((c06 >> 0) & 0xff) + |((c07 >> 0) & 0xff), 1, s_skb) + | BOX (((c13 >> 0) & 0xff) + |((c06 >> 8) & 0xff), 2, s_skb) + | BOX (((c20 >> 0) & 0xff) + |((c13 >> 8) & 0xff) + |((c06 >> 16) & 0xff), 3, s_skb); + + const u32x d00 = (d >> 0) & 0x00003c3f; + const u32x d07 = (d >> 7) & 0x00003f03; + const u32x d21 = (d >> 21) & 0x0000000f; + const u32x d22 = (d >> 22) & 0x00000030; + + u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb) + | BOX (((d07 >> 0) & 0xff) + |((d00 >> 8) & 0xff), 5, s_skb) + | BOX (((d07 >> 8) & 0xff), 6, s_skb) + | BOX (((d21 >> 0) & 0xff) + |((d22 >> 0) & 0xff), 7, s_skb); + + Kc[i] = ((t << 16) | (s & 0x0000ffff)); + Kd[i] = ((s >> 16) | (t & 0xffff0000)); + + Kc[i] = rotl32 (Kc[i], 2u); + Kd[i] = rotl32 (Kd[i], 2u); + } } -/* Fist0urs_end */ void AES256_ExpandKey (u32 *userkey, u32 *rek, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4) { @@ -215,54 +698,6 @@ void AES256_decrypt (const u32 *in, u32 *out, const u32 *rdk, SHM_TYPE u32 *s_td ^ rdk[59]; } -void AES256_encrypt (const u32 *in, u32 *out, const u32 *rek, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4) -{ - u32 t0 = in[0] ^ rek[0]; - u32 t1 = in[1] ^ rek[1]; - u32 t2 = in[2] ^ rek[2]; - u32 t3 = in[3] ^ rek[3]; - - #ifdef _unroll - #pragma unroll - #endif - for (int i = 4; i < 56; i += 4) - { - const uchar4 x0 = as_uchar4 (t0); - const uchar4 x1 = as_uchar4 (t1); - const uchar4 x2 = as_uchar4 (t2); - const uchar4 x3 = as_uchar4 (t3); - - t0 = s_te0[x0.s3] ^ s_te1[x1.s2] ^ s_te2[x2.s1] ^ s_te3[x3.s0] ^ rek[i + 0]; - t1 = s_te0[x1.s3] ^ s_te1[x2.s2] ^ s_te2[x3.s1] ^ s_te3[x0.s0] ^ rek[i + 1]; - t2 = s_te0[x2.s3] ^ s_te1[x3.s2] ^ s_te2[x0.s1] ^ s_te3[x1.s0] ^ rek[i + 2]; - t3 = s_te0[x3.s3] ^ s_te1[x0.s2] ^ s_te2[x1.s1] ^ s_te3[x2.s0] ^ rek[i + 3]; - } - - out[0] = (s_te4[(t0 >> 24) & 0xff] & 0xff000000) - ^ (s_te4[(t1 >> 16) & 0xff] & 0x00ff0000) - ^ (s_te4[(t2 >> 8) & 0xff] & 0x0000ff00) - ^ (s_te4[(t3 >> 0) & 0xff] & 0x000000ff) - ^ rek[56]; - - out[1] = (s_te4[(t1 >> 24) & 0xff] & 0xff000000) - ^ (s_te4[(t2 >> 16) & 0xff] & 0x00ff0000) - ^ (s_te4[(t3 >> 8) & 0xff] & 0x0000ff00) - ^ (s_te4[(t0 >> 0) & 0xff] & 0x000000ff) - ^ rek[57]; - - out[2] = (s_te4[(t2 >> 24) & 0xff] & 0xff000000) - ^ (s_te4[(t3 >> 16) & 0xff] & 0x00ff0000) - ^ (s_te4[(t0 >> 8) & 0xff] & 0x0000ff00) - ^ (s_te4[(t1 >> 0) & 0xff] & 0x000000ff) - ^ rek[58]; - - out[3] = (s_te4[(t3 >> 24) & 0xff] & 0xff000000) - ^ (s_te4[(t0 >> 16) & 0xff] & 0x00ff0000) - ^ (s_te4[(t1 >> 8) & 0xff] & 0x0000ff00) - ^ (s_te4[(t2 >> 0) & 0xff] & 0x000000ff) - ^ rek[59]; -} - #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" @@ -1107,27 +1542,7 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul * salt == SID */ - u32 salt_len = salt_bufs[salt_pos].salt_len; - - // u32 salt_buf0[4]; - // u32 salt_buf1[4]; - // u32 salt_buf2[4]; - // u32 salt_buf3[4]; - // u32 salt_buf4[4]; - // u32 salt_buf5[4]; - - // salt_buf0[0] = esalt_bufs[digests_offset].SID[0]; - // salt_buf0[1] = esalt_bufs[digests_offset].SID[1]; - // salt_buf0[2] = esalt_bufs[digests_offset].SID[2]; - // salt_buf0[3] = esalt_bufs[digests_offset].SID[3]; - // salt_buf1[0] = esalt_bufs[digests_offset].SID[4]; - // salt_buf1[1] = esalt_bufs[digests_offset].SID[5]; - // salt_buf1[2] = esalt_bufs[digests_offset].SID[6]; - // salt_buf1[3] = esalt_bufs[digests_offset].SID[7]; - // salt_buf2[0] = esalt_bufs[digests_offset].SID[8]; - // salt_buf2[1] = esalt_bufs[digests_offset].SID[9]; - // salt_buf2[2] = 0; - // salt_buf2[3] = 0; + const u32 salt_len = esalt_bufs[digests_offset].SID_len; u32 digest_context[5]; @@ -1264,9 +1679,15 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul u32 key[5]; hmac_sha1_run_S (w0, w1, w2, w3, ipad, opad, key); - /* this key is used as password for pbkdf2-hmac-* */ + esalt_bufs[digests_offset].userKey[0] = key[0]; + esalt_bufs[digests_offset].userKey[1] = key[1]; + esalt_bufs[digests_offset].userKey[2] = key[2]; + esalt_bufs[digests_offset].userKey[3] = key[3]; + esalt_bufs[digests_offset].userKey[4] = key[4]; + + /* if DPAPImk version 1, pbkdf-hmac-sha1 is used */ if (esalt_bufs[digests_offset].version == 1) { @@ -1309,13 +1730,13 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul * hmac1 */ - w0[0] = 0xdddddddd; - w0[1] = 0xdddddddd; - w0[2] = 0xdddddddd; - w0[3] = 0xdddddddd; - w1[0] = 0xdddddddd; - w1[1] = 1; - w1[2] = 0x80000000; + w0[0] = esalt_bufs[digests_offset].iv[0]; + w0[1] = esalt_bufs[digests_offset].iv[1]; + w0[2] = esalt_bufs[digests_offset].iv[2]; + w0[3] = esalt_bufs[digests_offset].iv[3]; + w1[0] = 1; + w1[1] = 0x80000000; + w1[2] = 0; w1[3] = 0; w2[0] = 0; w2[1] = 0; @@ -1324,7 +1745,7 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul w3[0] = 0; w3[1] = 0; w3[2] = 0; - w3[3] = (64 + 20 + 4) * 8; + w3[3] = (64 + 16 + 4) * 8; u32 digest[5]; @@ -1341,6 +1762,38 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].out[2] = digest[2]; tmps[gid].out[3] = digest[3]; tmps[gid].out[4] = digest[4]; + + /* We need bigger output ! */ + w0[0] = esalt_bufs[digests_offset].iv[0]; + w0[1] = esalt_bufs[digests_offset].iv[1]; + w0[2] = esalt_bufs[digests_offset].iv[2]; + w0[3] = esalt_bufs[digests_offset].iv[3]; + w1[0] = 2; + w1[1] = 0x80000000; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = (64 + 16 + 4) * 8; + + hmac_sha1_run_S (w0, w1, w2, w3, ipad, opad, digest); + + tmps[gid].dgst[5] = digest[0]; + tmps[gid].dgst[6] = digest[1]; + tmps[gid].dgst[7] = digest[2]; + tmps[gid].dgst[8] = digest[3]; + tmps[gid].dgst[9] = digest[4]; + + tmps[gid].out[5] = digest[0]; + tmps[gid].out[6] = digest[1]; + tmps[gid].out[7] = digest[2]; + tmps[gid].out[8] = digest[3]; + tmps[gid].out[9] = digest[4]; } /* if DPAPImk version 2, pbkdf-hmac-sha512 is used*/ else if (esalt_bufs[digests_offset].version == 2) @@ -1390,10 +1843,10 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].opad64[6] = opad64[6]; tmps[gid].opad64[7] = opad64[7]; - w0_x64[0] = hl32_to_64_S (0xdddddddd, 0xdddddddd); - w0_x64[1] = hl32_to_64_S (0xdddddddd, 0xdddddddd); - w0_x64[2] = hl32_to_64_S (0xdddddddd, 1); - w0_x64[3] = hl32_to_64_S (0x80000000, 0); + w0_x64[0] = hl32_to_64_S (esalt_bufs[digests_offset].iv[0], esalt_bufs[digests_offset].iv[1]); + w0_x64[1] = hl32_to_64_S (esalt_bufs[digests_offset].iv[2], esalt_bufs[digests_offset].iv[3]); + w0_x64[2] = hl32_to_64_S (1, 0x80000000); + w0_x64[3] = 0; w1_x64[0] = 0; w1_x64[1] = 0; w1_x64[2] = 0; @@ -1405,7 +1858,7 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul w3_x64[0] = 0; w3_x64[1] = 0; w3_x64[2] = 0; - w3_x64[3] = (128 + 20 + 4) * 8; + w3_x64[3] = (128 + 16 + 4) * 8; u64 dgst64[8]; @@ -1462,70 +1915,71 @@ __kernel void m15300_loop (__global pw_t *pws, __global const kernel_rule_t *rul * iter1 */ - u32x dgst[5]; - u32x out[5]; - - dgst[0] = packv (tmps, dgst, gid, 0); - dgst[1] = packv (tmps, dgst, gid, 1); - dgst[2] = packv (tmps, dgst, gid, 2); - dgst[3] = packv (tmps, dgst, gid, 3); - dgst[4] = packv (tmps, dgst, gid, 4); - - out[0] = packv (tmps, out, gid, 0); - out[1] = packv (tmps, out, gid, 1); - out[2] = packv (tmps, out, gid, 2); - out[3] = packv (tmps, out, gid, 3); - out[4] = packv (tmps, out, gid, 4); - - u8 hex[8] = {0}; - - for (u32 i = 0; i < loop_cnt; i++) + for (u32 i = 0; i < 8; i += 5) { - u32x w0[4]; - u32x w1[4]; - u32x w2[4]; - u32x w3[4]; + u32x dgst[5]; + u32x out[5]; + + dgst[0] = packv (tmps, dgst, gid, i + 0); + dgst[1] = packv (tmps, dgst, gid, i + 1); + dgst[2] = packv (tmps, dgst, gid, i + 2); + dgst[3] = packv (tmps, dgst, gid, i + 3); + dgst[4] = packv (tmps, dgst, gid, i + 4); + + out[0] = packv (tmps, out, gid, i + 0); + out[1] = packv (tmps, out, gid, i + 1); + out[2] = packv (tmps, out, gid, i + 2); + out[3] = packv (tmps, out, gid, i + 3); + out[4] = packv (tmps, out, gid, i + 4); - /* Microsoft PBKDF2 implementation. On purpose? - Misunderstanding of them? Dunno... - */ - w0[0] = out[0]; - w0[1] = out[1]; - w0[2] = out[2]; - w0[3] = out[3]; - w1[0] = out[4]; - w1[1] = 0x80000000; - w1[2] = 0; - w1[3] = 0; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = (64 + 20) * 8; - - hmac_sha1_run_V (w0, w1, w2, w3, ipad, opad, dgst); - - out[0] ^= dgst[0]; - out[1] ^= dgst[1]; - out[2] ^= dgst[2]; - out[3] ^= dgst[3]; - out[4] ^= dgst[4]; + for (u32 j = 0; j < loop_cnt; j++) + { + u32x w0[4]; + u32x w1[4]; + u32x w2[4]; + u32x w3[4]; + + /* Microsoft PBKDF2 implementation. On purpose? + Misunderstanding of them? Dunno... + */ + w0[0] = out[0]; + w0[1] = out[1]; + w0[2] = out[2]; + w0[3] = out[3]; + w1[0] = out[4]; + w1[1] = 0x80000000; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = (64 + 20) * 8; + + hmac_sha1_run_V (w0, w1, w2, w3, ipad, opad, dgst); + + out[0] ^= dgst[0]; + out[1] ^= dgst[1]; + out[2] ^= dgst[2]; + out[3] ^= dgst[3]; + out[4] ^= dgst[4]; + } + + unpackv (tmps, dgst, gid, i + 0, dgst[0]); + unpackv (tmps, dgst, gid, i + 1, dgst[1]); + unpackv (tmps, dgst, gid, i + 2, dgst[2]); + unpackv (tmps, dgst, gid, i + 3, dgst[3]); + unpackv (tmps, dgst, gid, i + 4, dgst[4]); + + unpackv (tmps, out, gid, i + 0, out[0]); + unpackv (tmps, out, gid, i + 1, out[1]); + unpackv (tmps, out, gid, i + 2, out[2]); + unpackv (tmps, out, gid, i + 3, out[3]); + unpackv (tmps, out, gid, i + 4, out[4]); } - - unpackv (tmps, dgst, gid, 0, dgst[0]); - unpackv (tmps, dgst, gid, 1, dgst[1]); - unpackv (tmps, dgst, gid, 2, dgst[2]); - unpackv (tmps, dgst, gid, 3, dgst[3]); - unpackv (tmps, dgst, gid, 4, dgst[4]); - - unpackv (tmps, out, gid, 0, out[0]); - unpackv (tmps, out, gid, 1, out[1]); - unpackv (tmps, out, gid, 2, out[2]); - unpackv (tmps, out, gid, 3, out[3]); - unpackv (tmps, out, gid, 4, out[4]); } else if (esalt_bufs[digests_offset].version == 2) { @@ -1632,16 +2086,511 @@ __kernel void m15300_loop (__global pw_t *pws, __global const kernel_rule_t *rul __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global dpapimk_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global dpapimk_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { - // u8 hex[16] = {0}; - // u64_to_hex_lower(r0, hex); - // printf("\n\nout[0]: %llu => ", r0); - // for (int i = 0; i < 16; i++) - // printf("%c", hex[i]); - // printf("\n"); - - // u64_to_hex_lower(r2, hex); - // printf("out[2]: %llu => ", r2); - // for (int i = 0; i < 16; i++) - // printf("%c", hex[i]); - // printf("\n"); + const u32 gid = get_global_id (0); + const u32 lid = get_local_id (0); + const u32 lsz = get_local_size (0); + + if (esalt_bufs[digests_offset].version == 1) + { + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + u32 ipad[5]; + u32 opad[5]; + + /** + * shared + */ + + __local u32 s_SPtrans[8][64]; + __local u32 s_skb[8][64]; + + for (u32 i = lid; i < 64; i += lsz) + { + s_SPtrans[0][i] = c_SPtrans[0][i]; + s_SPtrans[1][i] = c_SPtrans[1][i]; + s_SPtrans[2][i] = c_SPtrans[2][i]; + s_SPtrans[3][i] = c_SPtrans[3][i]; + s_SPtrans[4][i] = c_SPtrans[4][i]; + s_SPtrans[5][i] = c_SPtrans[5][i]; + s_SPtrans[6][i] = c_SPtrans[6][i]; + s_SPtrans[7][i] = c_SPtrans[7][i]; + + s_skb[0][i] = c_skb[0][i]; + s_skb[1][i] = c_skb[1][i]; + s_skb[2][i] = c_skb[2][i]; + s_skb[3][i] = c_skb[3][i]; + s_skb[4][i] = c_skb[4][i]; + s_skb[5][i] = c_skb[5][i]; + s_skb[6][i] = c_skb[6][i]; + s_skb[7][i] = c_skb[7][i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); + + if (gid >= gid_max) return; + + u32 key[6]; + + key[0] = swap32_S (tmps[gid].out[0]); + key[1] = swap32_S (tmps[gid].out[1]); + key[2] = swap32_S (tmps[gid].out[2]); + key[3] = swap32_S (tmps[gid].out[3]); + key[4] = swap32_S (tmps[gid].out[4]); + key[5] = swap32_S (tmps[gid].out[5]); + + u32x iv[2]; + + iv[0] = swap32_S (tmps[gid].out[6]); + iv[1] = swap32_S (tmps[gid].out[7]); + + u32 decrypted[26]; + + /* Construct 3DES keys */ + + const u32x a = (key[0]); + const u32x b = (key[1]); + + u32x Ka[16]; + u32x Kb[16]; + + _des_crypt_keysetup (a, b, Ka, Kb, s_skb); + + const u32x c = (key[2]); + const u32x d = (key[3]); + + u32x Kc[16]; + u32x Kd[16]; + + _des_crypt_keysetup (c, d, Kc, Kd, s_skb); + + const u32x e = (key[4]); + const u32x f = (key[5]); + + u32x Ke[16]; + u32x Kf[16]; + + _des_crypt_keysetup (e, f, Ke, Kf, s_skb); + + u32 contents_pos; + u32 contents_off; + u32 wx_off; + + for (wx_off = 0, contents_pos = 0, contents_off = 0; contents_pos < esalt_bufs[digests_offset].contents_len; wx_off += 2, contents_pos += 8, contents_off += 2) + { + /* First Pass */ + + u32x data[2]; + + data[0] = (u32x) swap32_S (esalt_bufs[digests_offset].contents[contents_off + 0]); + data[1] = (u32x) swap32_S (esalt_bufs[digests_offset].contents[contents_off + 1]); + + u32x p1[2]; + + _des_crypt_decrypt (p1, data, Ke, Kf, s_SPtrans); + + /* Second Pass */ + + u32x p2[2]; + + _des_crypt_encrypt (p2, p1, Kc, Kd, s_SPtrans); + + /* Third Pass */ + + u32x out[2]; + + _des_crypt_decrypt (out, p2, Ka, Kb, s_SPtrans); + + out[0] ^= iv[0]; + out[1] ^= iv[1]; + + decrypted[wx_off + 0] = out[0]; + decrypted[wx_off + 1] = out[1]; + + iv[0] = data[0]; + iv[1] = data[1]; + } + + u32 hmacSalt[4]; + u32 hmacTemoin[5]; + u32 lastKey[16]; + + hmacSalt[0] = swap32_S (decrypted[0]); + hmacSalt[1] = swap32_S (decrypted[1]); + hmacSalt[2] = swap32_S (decrypted[2]); + hmacSalt[3] = swap32_S (decrypted[3]); + + hmacTemoin[0] = swap32_S (decrypted[0 + 4]); + hmacTemoin[1] = swap32_S (decrypted[1 + 4]); + hmacTemoin[2] = swap32_S (decrypted[2 + 4]); + hmacTemoin[3] = swap32_S (decrypted[3 + 4]); + hmacTemoin[4] = swap32_S (decrypted[4 + 4]); + + for(int i = 0; i < 16; i++) + { + lastKey[i] = decrypted[i + 26 - 16]; + } + + w0[0] = esalt_bufs[digests_offset].userKey[0]; + w0[1] = esalt_bufs[digests_offset].userKey[1]; + w0[2] = esalt_bufs[digests_offset].userKey[2]; + w0[3] = esalt_bufs[digests_offset].userKey[3]; + w1[0] = esalt_bufs[digests_offset].userKey[4]; + w1[1] = 0; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + hmac_sha1_pad_S (w0, w1, w2, w3, ipad, opad); + + /** + * hmac1 + */ + + w0[0] = hmacSalt[0]; + w0[1] = hmacSalt[1]; + w0[2] = hmacSalt[2]; + w0[3] = hmacSalt[3]; + w1[0] = 0x80000000; + w1[1] = 0; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = (64 + 16) * 8; + + u32 digest[5]; + + hmac_sha1_run_S (w0, w1, w2, w3, ipad, opad, digest); + + w0[0] = digest[0]; + w0[1] = digest[1]; + w0[2] = digest[2]; + w0[3] = digest[3]; + w1[0] = digest[4]; + w1[1] = 0; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + hmac_sha1_pad_S (w0, w1, w2, w3, ipad, opad); + + /** + * hmac1 + */ + + w0[0] = swap32_S (lastKey[ 0]); + w0[1] = swap32_S (lastKey[ 1]); + w0[2] = swap32_S (lastKey[ 2]); + w0[3] = swap32_S (lastKey[ 3]); + w1[0] = swap32_S (lastKey[ 4]); + w1[1] = swap32_S (lastKey[ 5]); + w1[2] = swap32_S (lastKey[ 6]); + w1[3] = swap32_S (lastKey[ 7]); + w2[0] = swap32_S (lastKey[ 8]); + w2[1] = swap32_S (lastKey[ 9]); + w2[2] = swap32_S (lastKey[10]); + w2[3] = swap32_S (lastKey[11]); + w3[0] = swap32_S (lastKey[12]); + w3[1] = swap32_S (lastKey[13]); + w3[2] = swap32_S (lastKey[14]); + w3[3] = swap32_S (lastKey[15]); + + sha1_transform_S (w0, w1, w2, w3, ipad); + + w0[0] = 0x80000000; + w0[1] = 0; + w0[2] = 0; + w0[3] = 0; + w1[0] = 0; + w1[1] = 0; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = (64 + 16 * 4) * 8; + + hmac_sha1_run_S (w0, w1, w2, w3, ipad, opad, digest); + + #define il_pos 0 + + if ((hmacTemoin[0] == digest[0]) && (hmacTemoin[1] == digest[1]) + && (hmacTemoin[2] == digest[2]) && (hmacTemoin[3] == digest[3]) + && (hmacTemoin[4] == digest[4])) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + } + } + else if (esalt_bufs[digests_offset].version == 2) + { + /** + * aes shared + */ + + #ifdef REAL_SHM + + __local u32 s_td0[256]; + __local u32 s_td1[256]; + __local u32 s_td2[256]; + __local u32 s_td3[256]; + __local u32 s_td4[256]; + + __local u32 s_te0[256]; + __local u32 s_te1[256]; + __local u32 s_te2[256]; + __local u32 s_te3[256]; + __local u32 s_te4[256]; + + for (u32 i = lid; i < 256; i += lsz) + { + s_td0[i] = td0[i]; + s_td1[i] = td1[i]; + s_td2[i] = td2[i]; + s_td3[i] = td3[i]; + s_td4[i] = td4[i]; + + s_te0[i] = te0[i]; + s_te1[i] = te1[i]; + s_te2[i] = te2[i]; + s_te3[i] = te3[i]; + s_te4[i] = te4[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); + + #else + + __constant u32a *s_td0 = td0; + __constant u32a *s_td1 = td1; + __constant u32a *s_td2 = td2; + __constant u32a *s_td3 = td3; + __constant u32a *s_td4 = td4; + + __constant u32a *s_te0 = te0; + __constant u32a *s_te1 = te1; + __constant u32a *s_te2 = te2; + __constant u32a *s_te3 = te3; + __constant u32a *s_te4 = te4; + + #endif + + if (gid >= gid_max) return; + + /* Construct AES key */ + + u32 key[8]; + + key[0] = h32_from_64 (tmps[gid].out64[0]); + key[1] = l32_from_64 (tmps[gid].out64[0]); + key[2] = h32_from_64 (tmps[gid].out64[1]); + key[3] = l32_from_64 (tmps[gid].out64[1]); + key[4] = h32_from_64 (tmps[gid].out64[2]); + key[5] = l32_from_64 (tmps[gid].out64[2]); + key[6] = h32_from_64 (tmps[gid].out64[3]); + key[7] = l32_from_64 (tmps[gid].out64[3]); + + u32 iv[4]; + + iv[0] = h32_from_64 (tmps[gid].out64[4]); + iv[1] = l32_from_64 (tmps[gid].out64[4]); + iv[2] = h32_from_64 (tmps[gid].out64[5]); + iv[3] = l32_from_64 (tmps[gid].out64[5]); + + #define KEYLEN 60 + + u32 rek[KEYLEN]; + + AES256_ExpandKey (key, rek, s_te0, s_te1, s_te2, s_te3, s_te4); + + u32 rdk[KEYLEN]; + + #ifdef _unroll + #pragma unroll + #endif + for (u32 i = 0; i < KEYLEN; i++) rdk[i] = rek[i]; + + AES256_InvertKey (rdk, s_td0, s_td1, s_td2, s_td3, s_td4, s_te0, s_te1, s_te2, s_te3, s_te4); + + /* 144 bytes */ + u32 decrypted[36] = { 0 }; + + u32 contents_pos; + u32 contents_off; + u32 wx_off; + + for (wx_off = 0, contents_pos = 0, contents_off = 0; contents_pos < esalt_bufs[digests_offset].contents_len; wx_off += 4, contents_pos += 16, contents_off += 4) + { + u32 data[4]; + + data[0] = esalt_bufs[digests_offset].contents[contents_off + 0]; + data[1] = esalt_bufs[digests_offset].contents[contents_off + 1]; + data[2] = esalt_bufs[digests_offset].contents[contents_off + 2]; + data[3] = esalt_bufs[digests_offset].contents[contents_off + 3]; + + u32 out[4]; + + AES256_decrypt (data, out, rdk, s_td0, s_td1, s_td2, s_td3, s_td4); + + out[0] ^= iv[0]; + out[1] ^= iv[1]; + out[2] ^= iv[2]; + out[3] ^= iv[3]; + + decrypted[wx_off + 0] = out[0]; + decrypted[wx_off + 1] = out[1]; + decrypted[wx_off + 2] = out[2]; + decrypted[wx_off + 3] = out[3]; + + iv[0] = data[0]; + iv[1] = data[1]; + iv[2] = data[2]; + iv[3] = data[3]; + } + + u32 hmacSalt[4]; + u32 hmacTemoin[16]; + u32 lastKey[16]; + + hmacSalt[0] = decrypted[0]; + hmacSalt[1] = decrypted[1]; + hmacSalt[2] = decrypted[2]; + hmacSalt[3] = decrypted[3]; + + for(int i = 0; i < 16; i++) + { + hmacTemoin[i] = decrypted[i + 4]; + lastKey[i] = decrypted[i + 36 - 16]; + } + + u64 w0_x64[4]; + u64 w1_x64[4]; + u64 w2_x64[4]; + u64 w3_x64[4]; + + w0_x64[0] = hl32_to_64_S (esalt_bufs[digests_offset].userKey[0], esalt_bufs[digests_offset].userKey[1]); + w0_x64[1] = hl32_to_64_S (esalt_bufs[digests_offset].userKey[2], esalt_bufs[digests_offset].userKey[3]); + w0_x64[2] = hl32_to_64_S (esalt_bufs[digests_offset].userKey[4], 0); + w0_x64[3] = 0; + w1_x64[0] = 0; + w1_x64[1] = 0; + w1_x64[2] = 0; + w1_x64[3] = 0; + w2_x64[0] = 0; + w2_x64[1] = 0; + w2_x64[2] = 0; + w2_x64[3] = 0; + w3_x64[0] = 0; + w3_x64[1] = 0; + w3_x64[2] = 0; + w3_x64[3] = 0; + + u64 ipad64[8]; + u64 opad64[8]; + + hmac_sha512_pad_S (w0_x64, w1_x64, w2_x64, w3_x64, ipad64, opad64); + + w0_x64[0] = hl32_to_64_S (hmacSalt[0], hmacSalt[1]); + w0_x64[1] = hl32_to_64_S (hmacSalt[2], hmacSalt[3]); + w0_x64[2] = hl32_to_64_S (0x80000000, 0); + w0_x64[3] = 0; + w1_x64[0] = 0; + w1_x64[1] = 0; + w1_x64[2] = 0; + w1_x64[3] = 0; + w2_x64[0] = 0; + w2_x64[1] = 0; + w2_x64[2] = 0; + w2_x64[3] = 0; + w3_x64[0] = 0; + w3_x64[1] = 0; + w3_x64[2] = 0; + w3_x64[3] = (128 + 16) * 8; + + u64 dgst64[8]; + + hmac_sha512_run_S (w0_x64, w1_x64, w2_x64, w3_x64, ipad64, opad64, dgst64); + + u64 encKey[8]; + + encKey[0] = dgst64[0]; + encKey[1] = dgst64[1]; + encKey[2] = dgst64[2]; + encKey[3] = dgst64[3]; + encKey[4] = dgst64[4]; + encKey[5] = dgst64[5]; + encKey[6] = dgst64[6]; + encKey[7] = dgst64[7]; + + w0_x64[0] = encKey[0]; + w0_x64[1] = encKey[1]; + w0_x64[2] = encKey[2]; + w0_x64[3] = encKey[3]; + w1_x64[0] = encKey[4]; + w1_x64[1] = encKey[5]; + w1_x64[2] = encKey[6]; + w1_x64[3] = encKey[7]; + w2_x64[0] = 0; + w2_x64[1] = 0; + w2_x64[2] = 0; + w2_x64[3] = 0; + w3_x64[0] = 0; + w3_x64[1] = 0; + w3_x64[2] = 0; + w3_x64[3] = 0; + + hmac_sha512_pad_S (w0_x64, w1_x64, w2_x64, w3_x64, ipad64, opad64); + + w0_x64[0] = hl32_to_64_S (lastKey[ 0], lastKey[ 1]); + w0_x64[1] = hl32_to_64_S (lastKey[ 2], lastKey[ 3]); + w0_x64[2] = hl32_to_64_S (lastKey[ 4], lastKey[ 5]); + w0_x64[3] = hl32_to_64_S (lastKey[ 6], lastKey[ 7]); + w1_x64[0] = hl32_to_64_S (lastKey[ 8], lastKey[ 9]); + w1_x64[1] = hl32_to_64_S (lastKey[10], lastKey[11]); + w1_x64[2] = hl32_to_64_S (lastKey[12], lastKey[13]); + w1_x64[3] = hl32_to_64_S (lastKey[14], lastKey[15]); + w2_x64[0] = hl32_to_64_S (0x80000000, 0); + w2_x64[1] = 0; + w2_x64[2] = 0; + w2_x64[3] = 0; + w3_x64[0] = 0; + w3_x64[1] = 0; + w3_x64[2] = 0; + w3_x64[3] = (128 + 16 * 4) * 8; + + hmac_sha512_run_S (w0_x64, w1_x64, w2_x64, w3_x64, ipad64, opad64, dgst64); + + #define il_pos 0 + + if ((hmacTemoin[0] == h32_from_64 (dgst64[0])) && (hmacTemoin[1] == l32_from_64 (dgst64[0])) + && (hmacTemoin[14] == h32_from_64 (dgst64[7])) && (hmacTemoin[15] == l32_from_64 (dgst64[7]))) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + } + } } diff --git a/include/interface.h b/include/interface.h index 30bf2f8bd..4e84bd479 100644 --- a/include/interface.h +++ b/include/interface.h @@ -420,14 +420,16 @@ typedef struct psafe3 } psafe3_t; -/* Fist0urs */ typedef struct dpapimk { u32 version; u32 context; u8 SID_tmp[64]; - u32 SID[64]; + u32 SID[32]; + u32 SID_len; + u32 SID_offset; + u32 userKey[5]; u8 cipher_algo[16]; u8 hash_algo[16]; @@ -436,7 +438,6 @@ typedef struct dpapimk u32 contents[128]; } dpapimk_t; -/* Fist0urs_end */ typedef struct pdf14_tmp { @@ -828,14 +829,13 @@ typedef struct keepass_tmp } keepass_tmp_t; -/* Fist0urs */ typedef struct dpapimk_tmp { /* dedicated to hmac-sha1 */ u32 ipad[5]; u32 opad[5]; - u32 dgst[5]; - u32 out[5]; + u32 dgst[10]; + u32 out[10]; /* dedicated to hmac-sha512 */ u64 ipad64[8]; @@ -844,7 +844,6 @@ typedef struct dpapimk_tmp u64 out64[16]; } dpapimk_tmp_t; -/* Fist0urs_end */ typedef struct seven_zip_hook { @@ -1224,10 +1223,8 @@ typedef enum display_len DISPLAY_LEN_MAX_15100 = 6 + 6 + 1 + 8 + 1 + 28, DISPLAY_LEN_MIN_15200 = 1 + 10 + 1 + 2 + 1 + 1 + 1 + 1 + 1 + 64, DISPLAY_LEN_MAX_15200 = 1 + 10 + 1 + 2 + 1 + 8 + 1 + 5 + 1 + 20000, -/* Fist0urs */ DISPLAY_LEN_MIN_15300 = 1 + 7 + 1 + 1 + 1 + 1 + 1 + 10 + 1 + 4 + 1 + 4 + 1 + 1 + 1 + 32 + 1 + 3 + 128, DISPLAY_LEN_MAX_15300 = 1 + 7 + 1 + 1 + 1 + 1 + 1 + 100 + 1 + 6 + 1 + 6 + 1 + 10 + 1 + 32 + 1 + 4 + 1 + 512, -/* Fist0urs_end */ DISPLAY_LEN_MIN_99999 = 1, DISPLAY_LEN_MAX_99999 = 55, @@ -1555,9 +1552,7 @@ typedef enum kern_type KERN_TYPE_SKIP32 = 14900, KERN_TYPE_FILEZILLA_SERVER = 15000, KERN_TYPE_NETBSD_SHA1CRYPT = 15100, -/* Fist0urs */ KERN_TYPE_DPAPIMK = 15300, -/* Fist0urs_end */ KERN_TYPE_PLAINTEXT = 99999, } kern_type_t; @@ -1628,9 +1623,7 @@ typedef enum rounds_count ROUNDS_ITUNES102_BACKUP = 10000, ROUNDS_ATLASSIAN = 10000, ROUNDS_NETBSD_SHA1CRYPT = 20000, -/* Fist0urs */ - ROUNDS_DPAPIMK = 24000 - 1, // can be really different but fits jtr -test -/* Fist0urs_end */ + ROUNDS_DPAPIMK = 24000 - 1, // from 4000 to 24000 (possibly more) ROUNDS_STDOUT = 0 } rounds_count_t; @@ -1808,9 +1801,8 @@ int sha256b64s_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_bu int filezilla_server_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig); int netbsd_sha1crypt_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig); int atlassian_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig); -/* Fist0urs */ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig); -/* Fist0urs_end */ + /** * hook functions */ diff --git a/src/interface.c b/src/interface.c index 4cb0b5a64..82cad0aa7 100644 --- a/src/interface.c +++ b/src/interface.c @@ -239,9 +239,7 @@ static const char HT_14900[] = "Skip32 (PT = $salt, key = $pass)"; static const char HT_15000[] = "FileZilla Server >= 0.9.55"; static const char HT_15100[] = "Juniper/NetBSD sha1crypt"; static const char HT_15200[] = "Blockchain, My Wallet, V2"; -/* Fist0urs */ static const char HT_15300[] = "DPAPI masterkey file v1 and v2"; -/* Fist0urs_end */ static const char HT_99999[] = "Plaintext"; static const char HT_00011[] = "Joomla < 2.5.18"; @@ -319,9 +317,7 @@ static const char SIGNATURE_CRAM_MD5[] = "$cram_md5$"; static const char SIGNATURE_DCC2[] = "$DCC2$"; static const char SIGNATURE_DJANGOPBKDF2[] = "pbkdf2_sha256$"; static const char SIGNATURE_DJANGOSHA1[] = "sha1$"; -/* Fist0urs */ static const char SIGNATURE_DPAPIMK[] = "$DPAPImk$"; -/* Fist0urs_end */ static const char SIGNATURE_DRUPAL7[] = "$S$"; static const char SIGNATURE_ECRYPTFS[] = "$ecryptfs$"; static const char SIGNATURE_EPISERVER4[] = "$episerver$*1*"; @@ -2835,28 +2831,6 @@ int dcc2_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSE return (PARSER_OK); } -/* Fist0urs */ -int -pretty_print(char *message, char *data, int len) -{ - int g = 0; - for (int i = 0 ; i < len; i++) - { - if (g == 0) - { - printf("%s: ", message); - g++; - } - printf("%02x", data[i]); - } - printf("\n"); - - return 1; - -} -/* Fist0urs_end */ - -/* Fist0urs */ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig) { if ((input_len < DISPLAY_LEN_MIN_15300) || (input_len > DISPLAY_LEN_MAX_15300)) return (PARSER_GLOBAL_LENGTH); @@ -2884,15 +2858,10 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN u8 *contents_len_pos; u8 *contents_pos; - int dbg = 1; - version_pos = input_buf + 8 + 1; dpapimk->version = atoll ((const char *) version_pos); - if (dbg) - printf("version: %uld\n", dpapimk->version); - context_pos = (u8 *) strchr ((const char *) version_pos, '*'); if (context_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED); @@ -2901,9 +2870,6 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN dpapimk->context = atoll ((const char *) context_pos); - if (dbg) - printf("context: %uld\n", dpapimk->context); - SID_pos = (u8 *) strchr ((const char *) context_pos, '*'); if (SID_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED); @@ -2916,12 +2882,7 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN for (int i = 0; i < cipher_algo_pos - SID_pos; i++) dpapimk->SID_tmp[i] = SID_pos[i]; - /* Specific to DPAPI: needs trailing'\0' while computing hash */ dpapimk->SID_tmp[cipher_algo_pos - SID_pos] = '\0'; - dpapimk->SID_tmp[cipher_algo_pos - SID_pos + 1] = '\0'; - - if (dbg) - printf("SID: %s\n", dpapimk->SID_tmp); cipher_algo_pos++; @@ -2930,11 +2891,10 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN if (hash_algo_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED); for (int i = 0; i < hash_algo_pos - cipher_algo_pos; i++) + { dpapimk->cipher_algo[i] = cipher_algo_pos[i]; + } dpapimk->cipher_algo[hash_algo_pos - cipher_algo_pos] = '\0'; - - if (dbg) - printf("cipher_algo: %s\n", dpapimk->cipher_algo); hash_algo_pos++; @@ -2946,9 +2906,6 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN dpapimk->hash_algo[i] = hash_algo_pos[i]; dpapimk->hash_algo[rounds_pos - hash_algo_pos] = '\0'; - if (dbg) - printf("hash_algo: %s\n", dpapimk->hash_algo); - rounds_pos++; salt->salt_iter = (atoll ((const char *) rounds_pos)) - 1; @@ -2972,6 +2929,11 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN dpapimk->iv[2] = hex_to_u32 ((const u8 *) &iv_pos[16]); dpapimk->iv[3] = hex_to_u32 ((const u8 *) &iv_pos[24]); + dpapimk->iv[0] = byte_swap_32 (dpapimk->iv[0]); + dpapimk->iv[1] = byte_swap_32 (dpapimk->iv[1]); + dpapimk->iv[2] = byte_swap_32 (dpapimk->iv[2]); + dpapimk->iv[3] = byte_swap_32 (dpapimk->iv[3]); + contents_len_pos++; dpapimk->contents_len = (atoll ((const char *) contents_len_pos)); @@ -2991,9 +2953,6 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN if (end_line - contents_pos != dpapimk->contents_len) return (PARSER_SALT_LENGTH); - if (dbg) - printf("real_content_len: %uld\n", (u32)(end_line - contents_pos)); - for (u32 i = 0; i < dpapimk->contents_len / 4; i++) { dpapimk->contents[i] = hex_to_u32 ((const u8 *) &contents_pos[i * 8]); @@ -3008,41 +2967,31 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN u32 SID_len = cipher_algo_pos - 1 - SID_pos; - u8 *SID_unicode = (u8 *) hcmalloc ((SID_len + 1) * 2); + /* maximum size of SID supported */ + u8 *SID_unicode = (u8 *) hcmalloc (32 * 4); + memset (SID_unicode, 0, 32 * 4); - memset (SID_unicode, 0, (SID_len + 1) * 2); - - if (dbg) - printf("SID_len_before: %d\n", SID_len); - - /* Specific to DPAPI, SID + '\0' */ for (u32 i = 0; i < SID_len; i += 1) { - SID_unicode[i*2] = SID_pos[i]; + SID_unicode[i * 2] = SID_pos[i]; } - //SID_unicode[SID_len*2] = '\0'; - //SID_unicode[SID_len*2+1] = '\0'; + SID_unicode[(SID_len + 1) * 2] = 0x80; - salt->salt_len = (SID_len) * 2; + /* Specific to DPAPI: needs trailing '\0' while computing hash */ + dpapimk->SID_len = (SID_len + 1) * 2; - memcpy ((u8 *) dpapimk->SID, SID_unicode, salt->salt_len); + memcpy ((u8 *) dpapimk->SID, SID_unicode, 32 * 4); - u32 max = salt->salt_len / 4; - - if (salt->salt_len % 4) max++; - - for (u32 i = 0; i < max; i++) + for (u32 i = 0; i < 32; i++) { dpapimk->SID[i] = byte_swap_32 (dpapimk->SID[i]); } - dpapimk->SID[max] = 0x80000000; - + hcfree(SID_unicode); return (PARSER_OK); } -/* Fist0urs_end */ int wpa_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig) { @@ -18499,7 +18448,6 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le snprintf (out_buf, out_len - 1, "%s", hash_buf); } -/* Fist0urs */ else if (hash_mode == 15300) { dpapimk_t *dpapimks = (dpapimk_t *) esalts_buf; @@ -18509,26 +18457,34 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le u32 version = (u32) dpapimk->version; u32 context = (u32) dpapimk->context; u32 rounds = salt.salt_iter + 1; - u32 iv_len = 16; + u32 iv_len = 32; u32 contents_len = (u32) dpapimk->contents_len; char *ptr_SID = (char *) dpapimk->SID_tmp; char *ptr_cipher_algorithm = (char *) dpapimk->cipher_algo; char *ptr_hash_algorithm = (char *) dpapimk->hash_algo; - u8 *ptr_iv = (u8 *) dpapimk->iv; - u8 *ptr_contents = (u8 *) dpapimk->contents; + u32 *ptr_iv = (u32 *) dpapimk->iv; + u32 *ptr_contents = (u32 *) dpapimk->contents; - char data[16 * 2 + 1 + 288] = { 0 }; + u32 u32_iv[4]; + u8 iv[32 + 1]; - char *ptr_data = data; + for (u32 i = 0 ; i < iv_len / 8; i++) + { + u32_iv[i] = byte_swap_32 (ptr_iv[i]); + u32_to_hex_lower (u32_iv[i], iv + i * 8); + } + iv[32] = '\0'; - for (u32 i = 0; i < iv_len; i++, ptr_data += 2) - sprintf (ptr_data, "%02x", ptr_iv[i]); + u32 u32_contents[36]; + u8 contents[288 + 1]; - ptr_data++; - - for (u32 i = 0; i < contents_len; i++, ptr_data += 2) - sprintf (ptr_data, "%02x", ptr_contents[i]); + for (u32 i = 0 ; i < contents_len / 8; i++) + { + u32_contents[i] = byte_swap_32 (ptr_contents[i]); + u32_to_hex_lower (u32_contents[i], contents + i * 8); + } + contents[288] = '\0'; snprintf (out_buf, out_len - 1, "%s%d*%d*%s*%s*%s*%d*%s*%d*%s", SIGNATURE_DPAPIMK, @@ -18538,11 +18494,10 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le ptr_cipher_algorithm, ptr_hash_algorithm, rounds, - data, + iv, contents_len, - data + 33); + contents); } -/* Fist0urs_end */ else if (hash_mode == 99999) { char *ptr = (char *) digest_buf; @@ -22801,23 +22756,19 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->dgst_pos3 = 3; break; -/* Fist0urs */ case 15300: hashconfig->hash_type = HASH_TYPE_DPAPIMK; hashconfig->salt_type = SALT_TYPE_EMBEDDED; hashconfig->attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL; - hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_ST_UNICODE; + hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE; hashconfig->kern_type = KERN_TYPE_DPAPIMK; - hashconfig->dgst_size = DGST_SIZE_4_5; // because kernel uses _SHA1_ + hashconfig->dgst_size = DGST_SIZE_4_5; hashconfig->parse_func = dpapimk_parse_hash; - hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD; + hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE; hashconfig->dgst_pos0 = 0; hashconfig->dgst_pos1 = 1; hashconfig->dgst_pos2 = 2; hashconfig->dgst_pos3 = 3; break; -/* Fist0urs_end */ case 99999: hashconfig->hash_type = HASH_TYPE_PLAINTEXT; hashconfig->salt_type = SALT_TYPE_NONE; @@ -22954,9 +22905,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) case 14600: hashconfig->esalt_size = sizeof (luks_t); break; case 14700: hashconfig->esalt_size = sizeof (itunes_backup_t); break; case 14800: hashconfig->esalt_size = sizeof (itunes_backup_t); break; -/* Fist0urs */ case 15300: hashconfig->esalt_size = sizeof (dpapimk_t); break; -/* Fist0urs_end */ } // hook_salt_size @@ -23060,9 +23009,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) case 14800: hashconfig->tmp_size = sizeof (pbkdf2_sha256_tmp_t); break; case 15100: hashconfig->tmp_size = sizeof (pbkdf1_sha1_tmp_t); break; case 15200: hashconfig->tmp_size = sizeof (mywallet_tmp_t); break; -/* Fist0urs */ case 15300: hashconfig->tmp_size = sizeof (dpapimk_tmp_t); break; -/* Fist0urs_end */ }; // hook_size @@ -23467,10 +23414,8 @@ void hashconfig_benchmark_defaults (hashcat_ctx_t *hashcat_ctx, salt_t *salt, vo ((luks_t *) esalt)->cipher_type = HC_LUKS_CIPHER_TYPE_AES; ((luks_t *) esalt)->cipher_mode = HC_LUKS_CIPHER_MODE_XTS_PLAIN; break; -/* Fist0urs */ - case 15300: ((dpapimk_t *) esalt)->version = 2; + case 15300: ((dpapimk_t *) esalt)->version = 1; break; -/* Fist0urs_end */ } // special hook salt handling @@ -23665,10 +23610,8 @@ void hashconfig_benchmark_defaults (hashcat_ctx_t *hashcat_ctx, salt_t *salt, vo break; case 15200: salt->salt_iter = ROUNDS_MYWALLETV2; break; -/* Fist0urs */ case 15300: salt->salt_iter = ROUNDS_DPAPIMK; break; -/* Fist0urs_end */ } }