diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl index c4d69dbe7..a484814a6 100644 --- a/OpenCL/inc_types.cl +++ b/OpenCL/inc_types.cl @@ -1033,6 +1033,27 @@ typedef struct } keepass_t; +typedef struct +{ + u32 version; + u32 context; + + u32 SID[32]; + u32 SID_len; + u32 SID_offset; + + /* here only for possible + forward compatibiliy + */ + //char cipher_algo[16]; + //char hash_algo[16]; + + u32 iv[4]; + u32 contents_len; + u32 contents[128]; + +} dpapimk_t; + typedef struct { u32 digest[4]; @@ -1372,6 +1393,24 @@ typedef struct } keepass_tmp_t; +typedef struct +{ + /* dedicated to hmac-sha1 */ + u32 ipad[5]; + u32 opad[5]; + u32 dgst[10]; + u32 out[10]; + + u32 userKey[5]; + + /* dedicated to hmac-sha512 */ + u64 ipad64[8]; + u64 opad64[8]; + u64 dgst64[16]; + u64 out64[16]; + +} dpapimk_tmp_t; + typedef struct { u32 Kc[16]; diff --git a/OpenCL/m15300.cl b/OpenCL/m15300.cl new file mode 100644 index 000000000..184dcf447 --- /dev/null +++ b/OpenCL/m15300.cl @@ -0,0 +1,2634 @@ +/** + * Author......: See docs/credits.txt + * License.....: MIT + */ + +#define NEW_SIMD_CODE + +#include "inc_vendor.cl" +#include "inc_hash_constants.h" +#include "inc_hash_functions.cl" +#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" + +#define COMPARE_S "inc_comp_single.cl" +#define COMPARE_M "inc_comp_multi.cl" + +#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; \ +} + +#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; \ +} + +#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] = +{ + { + /* 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 + } +}; + +#define BOX(i,n,S) (S)[(n)][(i)] +#define BOX1(i,S) (S)[(i)] + +void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64]) +{ + u32 tt; + + u32 r = data[0]; + u32 l = data[1]; + + IP (r, l, tt); + + r = rotl32_S (r, 3u); + l = rotl32_S (l, 3u); + + #ifdef _unroll + #pragma unroll + #endif + for (u32 i = 0; i < 16; i += 2) + { + u32 u; + u32 t; + + u = Kc[i + 0] ^ r; + t = Kd[i + 0] ^ rotl32_S (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_S (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_S (l, 29u); + r = rotl32_S (r, 29u); + + FP (r, l, tt); + + iv[0] = l; + iv[1] = r; +} + +void _des_crypt_decrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64]) +{ + u32 tt; + + u32 r = data[0]; + u32 l = data[1]; + + IP (r, l, tt); + + r = rotl32_S (r, 3u); + l = rotl32_S (l, 3u); + + #ifdef _unroll + #pragma unroll + #endif + for (u32 i = 16; i > 0; i -= 2) + { + u32 u; + u32 t; + + u = Kc[i - 1] ^ r; + t = Kd[i - 1] ^ rotl32_S (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_S (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_S (l, 29u); + r = rotl32_S (r, 29u); + + FP (r, l, tt); + + iv[0] = l; + iv[1] = r; +} + +void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 (*s_skb)[64]) +{ + u32 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 u32 c00 = (c >> 0) & 0x0000003f; + const u32 c06 = (c >> 6) & 0x00383003; + const u32 c07 = (c >> 7) & 0x0000003c; + const u32 c13 = (c >> 13) & 0x0000060f; + const u32 c20 = (c >> 20) & 0x00000001; + + u32 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 u32 d00 = (d >> 0) & 0x00003c3f; + const u32 d07 = (d >> 7) & 0x00003f03; + const u32 d21 = (d >> 21) & 0x0000000f; + const u32 d22 = (d >> 22) & 0x00000030; + + u32 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_S (Kc[i], 2u); + Kd[i] = rotl32_S (Kd[i], 2u); + } +} + +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) +{ + rek[0] = userkey[0]; + rek[1] = userkey[1]; + rek[2] = userkey[2]; + rek[3] = userkey[3]; + rek[4] = userkey[4]; + rek[5] = userkey[5]; + rek[6] = userkey[6]; + rek[7] = userkey[7]; + + #ifdef _unroll + #pragma unroll + #endif + for (int i = 0, j = 0; i < 7; i += 1, j += 8) + { + const u32 temp1 = rek[j + 7]; + + rek[j + 8] = rek[j + 0] + ^ (s_te2[(temp1 >> 16) & 0xff] & 0xff000000) + ^ (s_te3[(temp1 >> 8) & 0xff] & 0x00ff0000) + ^ (s_te0[(temp1 >> 0) & 0xff] & 0x0000ff00) + ^ (s_te1[(temp1 >> 24) & 0xff] & 0x000000ff) + ^ rcon[i]; + rek[j + 9] = rek[j + 1] ^ rek[j + 8]; + rek[j + 10] = rek[j + 2] ^ rek[j + 9]; + rek[j + 11] = rek[j + 3] ^ rek[j + 10]; + + if (i == 6) continue; + + const u32 temp2 = rek[j + 11]; + + rek[j + 12] = rek[j + 4] + ^ (s_te2[(temp2 >> 24) & 0xff] & 0xff000000) + ^ (s_te3[(temp2 >> 16) & 0xff] & 0x00ff0000) + ^ (s_te0[(temp2 >> 8) & 0xff] & 0x0000ff00) + ^ (s_te1[(temp2 >> 0) & 0xff] & 0x000000ff); + rek[j + 13] = rek[j + 5] ^ rek[j + 12]; + rek[j + 14] = rek[j + 6] ^ rek[j + 13]; + rek[j + 15] = rek[j + 7] ^ rek[j + 14]; + } +} + +void AES256_InvertKey (u32 *rdk, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4, 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) +{ + #ifdef _unroll + #pragma unroll + #endif + for (u32 i = 0, j = 56; i < 28; i += 4, j -= 4) + { + u32 temp; + + temp = rdk[i + 0]; rdk[i + 0] = rdk[j + 0]; rdk[j + 0] = temp; + temp = rdk[i + 1]; rdk[i + 1] = rdk[j + 1]; rdk[j + 1] = temp; + temp = rdk[i + 2]; rdk[i + 2] = rdk[j + 2]; rdk[j + 2] = temp; + temp = rdk[i + 3]; rdk[i + 3] = rdk[j + 3]; rdk[j + 3] = temp; + } + + #ifdef _unroll + #pragma unroll + #endif + for (u32 i = 1, j = 4; i < 14; i += 1, j += 4) + { + rdk[j + 0] = + s_td0[s_te1[(rdk[j + 0] >> 24) & 0xff] & 0xff] ^ + s_td1[s_te1[(rdk[j + 0] >> 16) & 0xff] & 0xff] ^ + s_td2[s_te1[(rdk[j + 0] >> 8) & 0xff] & 0xff] ^ + s_td3[s_te1[(rdk[j + 0] >> 0) & 0xff] & 0xff]; + + rdk[j + 1] = + s_td0[s_te1[(rdk[j + 1] >> 24) & 0xff] & 0xff] ^ + s_td1[s_te1[(rdk[j + 1] >> 16) & 0xff] & 0xff] ^ + s_td2[s_te1[(rdk[j + 1] >> 8) & 0xff] & 0xff] ^ + s_td3[s_te1[(rdk[j + 1] >> 0) & 0xff] & 0xff]; + + rdk[j + 2] = + s_td0[s_te1[(rdk[j + 2] >> 24) & 0xff] & 0xff] ^ + s_td1[s_te1[(rdk[j + 2] >> 16) & 0xff] & 0xff] ^ + s_td2[s_te1[(rdk[j + 2] >> 8) & 0xff] & 0xff] ^ + s_td3[s_te1[(rdk[j + 2] >> 0) & 0xff] & 0xff]; + + rdk[j + 3] = + s_td0[s_te1[(rdk[j + 3] >> 24) & 0xff] & 0xff] ^ + s_td1[s_te1[(rdk[j + 3] >> 16) & 0xff] & 0xff] ^ + s_td2[s_te1[(rdk[j + 3] >> 8) & 0xff] & 0xff] ^ + s_td3[s_te1[(rdk[j + 3] >> 0) & 0xff] & 0xff]; + } +} + +void AES256_decrypt (const u32 *in, u32 *out, const u32 *rdk, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +{ + u32 t0 = in[0] ^ rdk[0]; + u32 t1 = in[1] ^ rdk[1]; + u32 t2 = in[2] ^ rdk[2]; + u32 t3 = in[3] ^ rdk[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_td0[x0.s3] ^ s_td1[x3.s2] ^ s_td2[x2.s1] ^ s_td3[x1.s0] ^ rdk[i + 0]; + t1 = s_td0[x1.s3] ^ s_td1[x0.s2] ^ s_td2[x3.s1] ^ s_td3[x2.s0] ^ rdk[i + 1]; + t2 = s_td0[x2.s3] ^ s_td1[x1.s2] ^ s_td2[x0.s1] ^ s_td3[x3.s0] ^ rdk[i + 2]; + t3 = s_td0[x3.s3] ^ s_td1[x2.s2] ^ s_td2[x1.s1] ^ s_td3[x0.s0] ^ rdk[i + 3]; + } + + out[0] = (s_td4[(t0 >> 24) & 0xff] & 0xff000000) + ^ (s_td4[(t3 >> 16) & 0xff] & 0x00ff0000) + ^ (s_td4[(t2 >> 8) & 0xff] & 0x0000ff00) + ^ (s_td4[(t1 >> 0) & 0xff] & 0x000000ff) + ^ rdk[56]; + + out[1] = (s_td4[(t1 >> 24) & 0xff] & 0xff000000) + ^ (s_td4[(t0 >> 16) & 0xff] & 0x00ff0000) + ^ (s_td4[(t3 >> 8) & 0xff] & 0x0000ff00) + ^ (s_td4[(t2 >> 0) & 0xff] & 0x000000ff) + ^ rdk[57]; + + out[2] = (s_td4[(t2 >> 24) & 0xff] & 0xff000000) + ^ (s_td4[(t1 >> 16) & 0xff] & 0x00ff0000) + ^ (s_td4[(t0 >> 8) & 0xff] & 0x0000ff00) + ^ (s_td4[(t3 >> 0) & 0xff] & 0x000000ff) + ^ rdk[58]; + + out[3] = (s_td4[(t3 >> 24) & 0xff] & 0xff000000) + ^ (s_td4[(t2 >> 16) & 0xff] & 0x00ff0000) + ^ (s_td4[(t1 >> 8) & 0xff] & 0x0000ff00) + ^ (s_td4[(t0 >> 0) & 0xff] & 0x000000ff) + ^ rdk[59]; +} + +void md4_transform_S (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4]) +{ + u32 a = digest[0]; + u32 b = digest[1]; + u32 c = digest[2]; + u32 d = digest[3]; + + MD4_STEP_S (MD4_Fo, a, b, c, d, w0[0], MD4C00, MD4S00); + MD4_STEP_S (MD4_Fo, d, a, b, c, w0[1], MD4C00, MD4S01); + MD4_STEP_S (MD4_Fo, c, d, a, b, w0[2], MD4C00, MD4S02); + MD4_STEP_S (MD4_Fo, b, c, d, a, w0[3], MD4C00, MD4S03); + MD4_STEP_S (MD4_Fo, a, b, c, d, w1[0], MD4C00, MD4S00); + MD4_STEP_S (MD4_Fo, d, a, b, c, w1[1], MD4C00, MD4S01); + MD4_STEP_S (MD4_Fo, c, d, a, b, w1[2], MD4C00, MD4S02); + MD4_STEP_S (MD4_Fo, b, c, d, a, w1[3], MD4C00, MD4S03); + MD4_STEP_S (MD4_Fo, a, b, c, d, w2[0], MD4C00, MD4S00); + MD4_STEP_S (MD4_Fo, d, a, b, c, w2[1], MD4C00, MD4S01); + MD4_STEP_S (MD4_Fo, c, d, a, b, w2[2], MD4C00, MD4S02); + MD4_STEP_S (MD4_Fo, b, c, d, a, w2[3], MD4C00, MD4S03); + MD4_STEP_S (MD4_Fo, a, b, c, d, w3[0], MD4C00, MD4S00); + MD4_STEP_S (MD4_Fo, d, a, b, c, w3[1], MD4C00, MD4S01); + MD4_STEP_S (MD4_Fo, c, d, a, b, w3[2], MD4C00, MD4S02); + MD4_STEP_S (MD4_Fo, b, c, d, a, w3[3], MD4C00, MD4S03); + + MD4_STEP_S (MD4_Go, a, b, c, d, w0[0], MD4C01, MD4S10); + MD4_STEP_S (MD4_Go, d, a, b, c, w1[0], MD4C01, MD4S11); + MD4_STEP_S (MD4_Go, c, d, a, b, w2[0], MD4C01, MD4S12); + MD4_STEP_S (MD4_Go, b, c, d, a, w3[0], MD4C01, MD4S13); + MD4_STEP_S (MD4_Go, a, b, c, d, w0[1], MD4C01, MD4S10); + MD4_STEP_S (MD4_Go, d, a, b, c, w1[1], MD4C01, MD4S11); + MD4_STEP_S (MD4_Go, c, d, a, b, w2[1], MD4C01, MD4S12); + MD4_STEP_S (MD4_Go, b, c, d, a, w3[1], MD4C01, MD4S13); + MD4_STEP_S (MD4_Go, a, b, c, d, w0[2], MD4C01, MD4S10); + MD4_STEP_S (MD4_Go, d, a, b, c, w1[2], MD4C01, MD4S11); + MD4_STEP_S (MD4_Go, c, d, a, b, w2[2], MD4C01, MD4S12); + MD4_STEP_S (MD4_Go, b, c, d, a, w3[2], MD4C01, MD4S13); + MD4_STEP_S (MD4_Go, a, b, c, d, w0[3], MD4C01, MD4S10); + MD4_STEP_S (MD4_Go, d, a, b, c, w1[3], MD4C01, MD4S11); + MD4_STEP_S (MD4_Go, c, d, a, b, w2[3], MD4C01, MD4S12); + MD4_STEP_S (MD4_Go, b, c, d, a, w3[3], MD4C01, MD4S13); + + MD4_STEP_S (MD4_H , a, b, c, d, w0[0], MD4C02, MD4S20); + MD4_STEP_S (MD4_H , d, a, b, c, w2[0], MD4C02, MD4S21); + MD4_STEP_S (MD4_H , c, d, a, b, w1[0], MD4C02, MD4S22); + MD4_STEP_S (MD4_H , b, c, d, a, w3[0], MD4C02, MD4S23); + MD4_STEP_S (MD4_H , a, b, c, d, w0[2], MD4C02, MD4S20); + MD4_STEP_S (MD4_H , d, a, b, c, w2[2], MD4C02, MD4S21); + MD4_STEP_S (MD4_H , c, d, a, b, w1[2], MD4C02, MD4S22); + MD4_STEP_S (MD4_H , b, c, d, a, w3[2], MD4C02, MD4S23); + MD4_STEP_S (MD4_H , a, b, c, d, w0[1], MD4C02, MD4S20); + MD4_STEP_S (MD4_H , d, a, b, c, w2[1], MD4C02, MD4S21); + MD4_STEP_S (MD4_H , c, d, a, b, w1[1], MD4C02, MD4S22); + MD4_STEP_S (MD4_H , b, c, d, a, w3[1], MD4C02, MD4S23); + MD4_STEP_S (MD4_H , a, b, c, d, w0[3], MD4C02, MD4S20); + MD4_STEP_S (MD4_H , d, a, b, c, w2[3], MD4C02, MD4S21); + MD4_STEP_S (MD4_H , c, d, a, b, w1[3], MD4C02, MD4S22); + MD4_STEP_S (MD4_H , b, c, d, a, w3[3], MD4C02, MD4S23); + + digest[0] += a; + digest[1] += b; + digest[2] += c; + digest[3] += d; +} + +void sha1_transform_S (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[5]) +{ + u32 A = digest[0]; + u32 B = digest[1]; + u32 C = digest[2]; + u32 D = digest[3]; + u32 E = digest[4]; + + u32 w0_t = w0[0]; + u32 w1_t = w0[1]; + u32 w2_t = w0[2]; + u32 w3_t = w0[3]; + u32 w4_t = w1[0]; + u32 w5_t = w1[1]; + u32 w6_t = w1[2]; + u32 w7_t = w1[3]; + u32 w8_t = w2[0]; + u32 w9_t = w2[1]; + u32 wa_t = w2[2]; + u32 wb_t = w2[3]; + u32 wc_t = w3[0]; + u32 wd_t = w3[1]; + u32 we_t = w3[2]; + u32 wf_t = w3[3]; + + #undef K + #define K SHA1C00 + + SHA1_STEP_S (SHA1_F0o, A, B, C, D, E, w0_t); + SHA1_STEP_S (SHA1_F0o, E, A, B, C, D, w1_t); + SHA1_STEP_S (SHA1_F0o, D, E, A, B, C, w2_t); + SHA1_STEP_S (SHA1_F0o, C, D, E, A, B, w3_t); + SHA1_STEP_S (SHA1_F0o, B, C, D, E, A, w4_t); + SHA1_STEP_S (SHA1_F0o, A, B, C, D, E, w5_t); + SHA1_STEP_S (SHA1_F0o, E, A, B, C, D, w6_t); + SHA1_STEP_S (SHA1_F0o, D, E, A, B, C, w7_t); + SHA1_STEP_S (SHA1_F0o, C, D, E, A, B, w8_t); + SHA1_STEP_S (SHA1_F0o, B, C, D, E, A, w9_t); + SHA1_STEP_S (SHA1_F0o, A, B, C, D, E, wa_t); + SHA1_STEP_S (SHA1_F0o, E, A, B, C, D, wb_t); + SHA1_STEP_S (SHA1_F0o, D, E, A, B, C, wc_t); + SHA1_STEP_S (SHA1_F0o, C, D, E, A, B, wd_t); + SHA1_STEP_S (SHA1_F0o, B, C, D, E, A, we_t); + SHA1_STEP_S (SHA1_F0o, A, B, C, D, E, wf_t); + w0_t = rotl32_S ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP_S (SHA1_F0o, E, A, B, C, D, w0_t); + w1_t = rotl32_S ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP_S (SHA1_F0o, D, E, A, B, C, w1_t); + w2_t = rotl32_S ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP_S (SHA1_F0o, C, D, E, A, B, w2_t); + w3_t = rotl32_S ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP_S (SHA1_F0o, B, C, D, E, A, w3_t); + + #undef K + #define K SHA1C01 + + w4_t = rotl32_S ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP_S (SHA1_F1, A, B, C, D, E, w4_t); + w5_t = rotl32_S ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP_S (SHA1_F1, E, A, B, C, D, w5_t); + w6_t = rotl32_S ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP_S (SHA1_F1, D, E, A, B, C, w6_t); + w7_t = rotl32_S ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP_S (SHA1_F1, C, D, E, A, B, w7_t); + w8_t = rotl32_S ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP_S (SHA1_F1, B, C, D, E, A, w8_t); + w9_t = rotl32_S ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP_S (SHA1_F1, A, B, C, D, E, w9_t); + wa_t = rotl32_S ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP_S (SHA1_F1, E, A, B, C, D, wa_t); + wb_t = rotl32_S ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP_S (SHA1_F1, D, E, A, B, C, wb_t); + wc_t = rotl32_S ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP_S (SHA1_F1, C, D, E, A, B, wc_t); + wd_t = rotl32_S ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP_S (SHA1_F1, B, C, D, E, A, wd_t); + we_t = rotl32_S ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP_S (SHA1_F1, A, B, C, D, E, we_t); + wf_t = rotl32_S ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP_S (SHA1_F1, E, A, B, C, D, wf_t); + w0_t = rotl32_S ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP_S (SHA1_F1, D, E, A, B, C, w0_t); + w1_t = rotl32_S ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP_S (SHA1_F1, C, D, E, A, B, w1_t); + w2_t = rotl32_S ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP_S (SHA1_F1, B, C, D, E, A, w2_t); + w3_t = rotl32_S ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP_S (SHA1_F1, A, B, C, D, E, w3_t); + w4_t = rotl32_S ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP_S (SHA1_F1, E, A, B, C, D, w4_t); + w5_t = rotl32_S ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP_S (SHA1_F1, D, E, A, B, C, w5_t); + w6_t = rotl32_S ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP_S (SHA1_F1, C, D, E, A, B, w6_t); + w7_t = rotl32_S ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP_S (SHA1_F1, B, C, D, E, A, w7_t); + + #undef K + #define K SHA1C02 + + w8_t = rotl32_S ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP_S (SHA1_F2o, A, B, C, D, E, w8_t); + w9_t = rotl32_S ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP_S (SHA1_F2o, E, A, B, C, D, w9_t); + wa_t = rotl32_S ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP_S (SHA1_F2o, D, E, A, B, C, wa_t); + wb_t = rotl32_S ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP_S (SHA1_F2o, C, D, E, A, B, wb_t); + wc_t = rotl32_S ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP_S (SHA1_F2o, B, C, D, E, A, wc_t); + wd_t = rotl32_S ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP_S (SHA1_F2o, A, B, C, D, E, wd_t); + we_t = rotl32_S ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP_S (SHA1_F2o, E, A, B, C, D, we_t); + wf_t = rotl32_S ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP_S (SHA1_F2o, D, E, A, B, C, wf_t); + w0_t = rotl32_S ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP_S (SHA1_F2o, C, D, E, A, B, w0_t); + w1_t = rotl32_S ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP_S (SHA1_F2o, B, C, D, E, A, w1_t); + w2_t = rotl32_S ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP_S (SHA1_F2o, A, B, C, D, E, w2_t); + w3_t = rotl32_S ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP_S (SHA1_F2o, E, A, B, C, D, w3_t); + w4_t = rotl32_S ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP_S (SHA1_F2o, D, E, A, B, C, w4_t); + w5_t = rotl32_S ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP_S (SHA1_F2o, C, D, E, A, B, w5_t); + w6_t = rotl32_S ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP_S (SHA1_F2o, B, C, D, E, A, w6_t); + w7_t = rotl32_S ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP_S (SHA1_F2o, A, B, C, D, E, w7_t); + w8_t = rotl32_S ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP_S (SHA1_F2o, E, A, B, C, D, w8_t); + w9_t = rotl32_S ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP_S (SHA1_F2o, D, E, A, B, C, w9_t); + wa_t = rotl32_S ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP_S (SHA1_F2o, C, D, E, A, B, wa_t); + wb_t = rotl32_S ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP_S (SHA1_F2o, B, C, D, E, A, wb_t); + + #undef K + #define K SHA1C03 + + wc_t = rotl32_S ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP_S (SHA1_F1, A, B, C, D, E, wc_t); + wd_t = rotl32_S ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP_S (SHA1_F1, E, A, B, C, D, wd_t); + we_t = rotl32_S ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP_S (SHA1_F1, D, E, A, B, C, we_t); + wf_t = rotl32_S ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP_S (SHA1_F1, C, D, E, A, B, wf_t); + w0_t = rotl32_S ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP_S (SHA1_F1, B, C, D, E, A, w0_t); + w1_t = rotl32_S ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP_S (SHA1_F1, A, B, C, D, E, w1_t); + w2_t = rotl32_S ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP_S (SHA1_F1, E, A, B, C, D, w2_t); + w3_t = rotl32_S ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP_S (SHA1_F1, D, E, A, B, C, w3_t); + w4_t = rotl32_S ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP_S (SHA1_F1, C, D, E, A, B, w4_t); + w5_t = rotl32_S ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP_S (SHA1_F1, B, C, D, E, A, w5_t); + w6_t = rotl32_S ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP_S (SHA1_F1, A, B, C, D, E, w6_t); + w7_t = rotl32_S ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP_S (SHA1_F1, E, A, B, C, D, w7_t); + w8_t = rotl32_S ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP_S (SHA1_F1, D, E, A, B, C, w8_t); + w9_t = rotl32_S ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP_S (SHA1_F1, C, D, E, A, B, w9_t); + wa_t = rotl32_S ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP_S (SHA1_F1, B, C, D, E, A, wa_t); + wb_t = rotl32_S ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP_S (SHA1_F1, A, B, C, D, E, wb_t); + wc_t = rotl32_S ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP_S (SHA1_F1, E, A, B, C, D, wc_t); + wd_t = rotl32_S ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP_S (SHA1_F1, D, E, A, B, C, wd_t); + we_t = rotl32_S ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP_S (SHA1_F1, C, D, E, A, B, we_t); + wf_t = rotl32_S ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP_S (SHA1_F1, B, C, D, E, A, wf_t); + + digest[0] += A; + digest[1] += B; + digest[2] += C; + digest[3] += D; + digest[4] += E; +} + +void hmac_sha1_pad_S (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[5], u32 opad[5]) +{ + w0[0] = w0[0] ^ 0x36363636; + w0[1] = w0[1] ^ 0x36363636; + w0[2] = w0[2] ^ 0x36363636; + w0[3] = w0[3] ^ 0x36363636; + w1[0] = w1[0] ^ 0x36363636; + w1[1] = w1[1] ^ 0x36363636; + w1[2] = w1[2] ^ 0x36363636; + w1[3] = w1[3] ^ 0x36363636; + w2[0] = w2[0] ^ 0x36363636; + w2[1] = w2[1] ^ 0x36363636; + w2[2] = w2[2] ^ 0x36363636; + w2[3] = w2[3] ^ 0x36363636; + w3[0] = w3[0] ^ 0x36363636; + w3[1] = w3[1] ^ 0x36363636; + w3[2] = w3[2] ^ 0x36363636; + w3[3] = w3[3] ^ 0x36363636; + + ipad[0] = SHA1M_A; + ipad[1] = SHA1M_B; + ipad[2] = SHA1M_C; + ipad[3] = SHA1M_D; + ipad[4] = SHA1M_E; + + sha1_transform_S (w0, w1, w2, w3, ipad); + + w0[0] = w0[0] ^ 0x6a6a6a6a; + w0[1] = w0[1] ^ 0x6a6a6a6a; + w0[2] = w0[2] ^ 0x6a6a6a6a; + w0[3] = w0[3] ^ 0x6a6a6a6a; + w1[0] = w1[0] ^ 0x6a6a6a6a; + w1[1] = w1[1] ^ 0x6a6a6a6a; + w1[2] = w1[2] ^ 0x6a6a6a6a; + w1[3] = w1[3] ^ 0x6a6a6a6a; + w2[0] = w2[0] ^ 0x6a6a6a6a; + w2[1] = w2[1] ^ 0x6a6a6a6a; + w2[2] = w2[2] ^ 0x6a6a6a6a; + w2[3] = w2[3] ^ 0x6a6a6a6a; + w3[0] = w3[0] ^ 0x6a6a6a6a; + w3[1] = w3[1] ^ 0x6a6a6a6a; + w3[2] = w3[2] ^ 0x6a6a6a6a; + w3[3] = w3[3] ^ 0x6a6a6a6a; + + opad[0] = SHA1M_A; + opad[1] = SHA1M_B; + opad[2] = SHA1M_C; + opad[3] = SHA1M_D; + opad[4] = SHA1M_E; + + sha1_transform_S (w0, w1, w2, w3, opad); +} + +void hmac_sha1_run_S (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[5], u32 opad[5], u32 digest[5]) +{ + digest[0] = ipad[0]; + digest[1] = ipad[1]; + digest[2] = ipad[2]; + digest[3] = ipad[3]; + digest[4] = ipad[4]; + + sha1_transform_S (w0, w1, w2, w3, digest); + + w0[0] = digest[0]; + w0[1] = digest[1]; + w0[2] = digest[2]; + w0[3] = digest[3]; + w1[0] = digest[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; + + digest[0] = opad[0]; + digest[1] = opad[1]; + digest[2] = opad[2]; + digest[3] = opad[3]; + digest[4] = opad[4]; + + sha1_transform_S (w0, w1, w2, w3, digest); +} + +void sha1_transform_V (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5]) +{ + u32x A = digest[0]; + u32x B = digest[1]; + u32x C = digest[2]; + u32x D = digest[3]; + u32x E = digest[4]; + + u32x w0_t = w0[0]; + u32x w1_t = w0[1]; + u32x w2_t = w0[2]; + u32x w3_t = w0[3]; + u32x w4_t = w1[0]; + u32x w5_t = w1[1]; + u32x w6_t = w1[2]; + u32x w7_t = w1[3]; + u32x w8_t = w2[0]; + u32x w9_t = w2[1]; + u32x wa_t = w2[2]; + u32x wb_t = w2[3]; + u32x wc_t = w3[0]; + u32x wd_t = w3[1]; + u32x we_t = w3[2]; + u32x wf_t = w3[3]; + + #undef K + #define K SHA1C00 + + SHA1_STEP (SHA1_F0o, A, B, C, D, E, w0_t); + SHA1_STEP (SHA1_F0o, E, A, B, C, D, w1_t); + SHA1_STEP (SHA1_F0o, D, E, A, B, C, w2_t); + SHA1_STEP (SHA1_F0o, C, D, E, A, B, w3_t); + SHA1_STEP (SHA1_F0o, B, C, D, E, A, w4_t); + SHA1_STEP (SHA1_F0o, A, B, C, D, E, w5_t); + SHA1_STEP (SHA1_F0o, E, A, B, C, D, w6_t); + SHA1_STEP (SHA1_F0o, D, E, A, B, C, w7_t); + SHA1_STEP (SHA1_F0o, C, D, E, A, B, w8_t); + SHA1_STEP (SHA1_F0o, B, C, D, E, A, w9_t); + SHA1_STEP (SHA1_F0o, A, B, C, D, E, wa_t); + SHA1_STEP (SHA1_F0o, E, A, B, C, D, wb_t); + SHA1_STEP (SHA1_F0o, D, E, A, B, C, wc_t); + SHA1_STEP (SHA1_F0o, C, D, E, A, B, wd_t); + SHA1_STEP (SHA1_F0o, B, C, D, E, A, we_t); + SHA1_STEP (SHA1_F0o, A, B, C, D, E, wf_t); + w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, E, A, B, C, D, w0_t); + w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, D, E, A, B, C, w1_t); + w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, C, D, E, A, B, w2_t); + w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, B, C, D, E, A, w3_t); + + #undef K + #define K SHA1C01 + + w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w4_t); + w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w5_t); + w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w6_t); + w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w7_t); + w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w8_t); + w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w9_t); + wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wa_t); + wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wb_t); + wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wc_t); + wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wd_t); + we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, we_t); + wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wf_t); + w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w0_t); + w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w1_t); + w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w2_t); + w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w3_t); + w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w4_t); + w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w5_t); + w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w6_t); + w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w7_t); + + #undef K + #define K SHA1C02 + + w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w8_t); + w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w9_t); + wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wa_t); + wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wb_t); + wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wc_t); + wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, wd_t); + we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, we_t); + wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wf_t); + w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w0_t); + w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w1_t); + w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w2_t); + w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w3_t); + w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w4_t); + w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w5_t); + w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w6_t); + w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w7_t); + w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w8_t); + w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w9_t); + wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wa_t); + wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wb_t); + + #undef K + #define K SHA1C03 + + wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wc_t); + wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wd_t); + we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, we_t); + wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wf_t); + w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w0_t); + w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w1_t); + w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w2_t); + w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w3_t); + w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w4_t); + w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w5_t); + w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w6_t); + w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w7_t); + w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w8_t); + w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w9_t); + wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wa_t); + wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wb_t); + wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wc_t); + wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wd_t); + we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, we_t); + wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wf_t); + + digest[0] += A; + digest[1] += B; + digest[2] += C; + digest[3] += D; + digest[4] += E; +} + +void hmac_sha1_pad_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[5], u32x opad[5]) +{ + w0[0] = w0[0] ^ 0x36363636; + w0[1] = w0[1] ^ 0x36363636; + w0[2] = w0[2] ^ 0x36363636; + w0[3] = w0[3] ^ 0x36363636; + w1[0] = w1[0] ^ 0x36363636; + w1[1] = w1[1] ^ 0x36363636; + w1[2] = w1[2] ^ 0x36363636; + w1[3] = w1[3] ^ 0x36363636; + w2[0] = w2[0] ^ 0x36363636; + w2[1] = w2[1] ^ 0x36363636; + w2[2] = w2[2] ^ 0x36363636; + w2[3] = w2[3] ^ 0x36363636; + w3[0] = w3[0] ^ 0x36363636; + w3[1] = w3[1] ^ 0x36363636; + w3[2] = w3[2] ^ 0x36363636; + w3[3] = w3[3] ^ 0x36363636; + + ipad[0] = SHA1M_A; + ipad[1] = SHA1M_B; + ipad[2] = SHA1M_C; + ipad[3] = SHA1M_D; + ipad[4] = SHA1M_E; + + sha1_transform_V (w0, w1, w2, w3, ipad); + + w0[0] = w0[0] ^ 0x6a6a6a6a; + w0[1] = w0[1] ^ 0x6a6a6a6a; + w0[2] = w0[2] ^ 0x6a6a6a6a; + w0[3] = w0[3] ^ 0x6a6a6a6a; + w1[0] = w1[0] ^ 0x6a6a6a6a; + w1[1] = w1[1] ^ 0x6a6a6a6a; + w1[2] = w1[2] ^ 0x6a6a6a6a; + w1[3] = w1[3] ^ 0x6a6a6a6a; + w2[0] = w2[0] ^ 0x6a6a6a6a; + w2[1] = w2[1] ^ 0x6a6a6a6a; + w2[2] = w2[2] ^ 0x6a6a6a6a; + w2[3] = w2[3] ^ 0x6a6a6a6a; + w3[0] = w3[0] ^ 0x6a6a6a6a; + w3[1] = w3[1] ^ 0x6a6a6a6a; + w3[2] = w3[2] ^ 0x6a6a6a6a; + w3[3] = w3[3] ^ 0x6a6a6a6a; + + opad[0] = SHA1M_A; + opad[1] = SHA1M_B; + opad[2] = SHA1M_C; + opad[3] = SHA1M_D; + opad[4] = SHA1M_E; + + sha1_transform_V (w0, w1, w2, w3, opad); +} + +void hmac_sha1_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[5], u32x opad[5], u32x digest[5]) +{ + digest[0] = ipad[0]; + digest[1] = ipad[1]; + digest[2] = ipad[2]; + digest[3] = ipad[3]; + digest[4] = ipad[4]; + + sha1_transform_V (w0, w1, w2, w3, digest); + + w0[0] = digest[0]; + w0[1] = digest[1]; + w0[2] = digest[2]; + w0[3] = digest[3]; + w1[0] = digest[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; + + digest[0] = opad[0]; + digest[1] = opad[1]; + digest[2] = opad[2]; + digest[3] = opad[3]; + digest[4] = opad[4]; + + sha1_transform_V (w0, w1, w2, w3, digest); +} + +__constant u64a k_sha512[80] = +{ + SHA512C00, SHA512C01, SHA512C02, SHA512C03, + SHA512C04, SHA512C05, SHA512C06, SHA512C07, + SHA512C08, SHA512C09, SHA512C0a, SHA512C0b, + SHA512C0c, SHA512C0d, SHA512C0e, SHA512C0f, + SHA512C10, SHA512C11, SHA512C12, SHA512C13, + SHA512C14, SHA512C15, SHA512C16, SHA512C17, + SHA512C18, SHA512C19, SHA512C1a, SHA512C1b, + SHA512C1c, SHA512C1d, SHA512C1e, SHA512C1f, + SHA512C20, SHA512C21, SHA512C22, SHA512C23, + SHA512C24, SHA512C25, SHA512C26, SHA512C27, + SHA512C28, SHA512C29, SHA512C2a, SHA512C2b, + SHA512C2c, SHA512C2d, SHA512C2e, SHA512C2f, + SHA512C30, SHA512C31, SHA512C32, SHA512C33, + SHA512C34, SHA512C35, SHA512C36, SHA512C37, + SHA512C38, SHA512C39, SHA512C3a, SHA512C3b, + SHA512C3c, SHA512C3d, SHA512C3e, SHA512C3f, + SHA512C40, SHA512C41, SHA512C42, SHA512C43, + SHA512C44, SHA512C45, SHA512C46, SHA512C47, + SHA512C48, SHA512C49, SHA512C4a, SHA512C4b, + SHA512C4c, SHA512C4d, SHA512C4e, SHA512C4f, +}; + +void sha512_transform_S (const u64 w0[4], const u64 w1[4], const u64 w2[4], const u64 w3[4], u64 digest[8]) +{ + u64 a = digest[0]; + u64 b = digest[1]; + u64 c = digest[2]; + u64 d = digest[3]; + u64 e = digest[4]; + u64 f = digest[5]; + u64 g = digest[6]; + u64 h = digest[7]; + + u64 w0_t = w0[0]; + u64 w1_t = w0[1]; + u64 w2_t = w0[2]; + u64 w3_t = w0[3]; + u64 w4_t = w1[0]; + u64 w5_t = w1[1]; + u64 w6_t = w1[2]; + u64 w7_t = w1[3]; + u64 w8_t = w2[0]; + u64 w9_t = w2[1]; + u64 wa_t = w2[2]; + u64 wb_t = w2[3]; + u64 wc_t = w3[0]; + u64 wd_t = w3[1]; + u64 we_t = w3[2]; + u64 wf_t = w3[3]; + + #define ROUND_EXPAND_S() \ + { \ + w0_t = SHA512_EXPAND_S (we_t, w9_t, w1_t, w0_t); \ + w1_t = SHA512_EXPAND_S (wf_t, wa_t, w2_t, w1_t); \ + w2_t = SHA512_EXPAND_S (w0_t, wb_t, w3_t, w2_t); \ + w3_t = SHA512_EXPAND_S (w1_t, wc_t, w4_t, w3_t); \ + w4_t = SHA512_EXPAND_S (w2_t, wd_t, w5_t, w4_t); \ + w5_t = SHA512_EXPAND_S (w3_t, we_t, w6_t, w5_t); \ + w6_t = SHA512_EXPAND_S (w4_t, wf_t, w7_t, w6_t); \ + w7_t = SHA512_EXPAND_S (w5_t, w0_t, w8_t, w7_t); \ + w8_t = SHA512_EXPAND_S (w6_t, w1_t, w9_t, w8_t); \ + w9_t = SHA512_EXPAND_S (w7_t, w2_t, wa_t, w9_t); \ + wa_t = SHA512_EXPAND_S (w8_t, w3_t, wb_t, wa_t); \ + wb_t = SHA512_EXPAND_S (w9_t, w4_t, wc_t, wb_t); \ + wc_t = SHA512_EXPAND_S (wa_t, w5_t, wd_t, wc_t); \ + wd_t = SHA512_EXPAND_S (wb_t, w6_t, we_t, wd_t); \ + we_t = SHA512_EXPAND_S (wc_t, w7_t, wf_t, we_t); \ + wf_t = SHA512_EXPAND_S (wd_t, w8_t, w0_t, wf_t); \ + } + + #define ROUND_STEP_S(i) \ + { \ + SHA512_STEP_S (SHA512_F0o, SHA512_F1o, a, b, c, d, e, f, g, h, w0_t, k_sha512[i + 0]); \ + SHA512_STEP_S (SHA512_F0o, SHA512_F1o, h, a, b, c, d, e, f, g, w1_t, k_sha512[i + 1]); \ + SHA512_STEP_S (SHA512_F0o, SHA512_F1o, g, h, a, b, c, d, e, f, w2_t, k_sha512[i + 2]); \ + SHA512_STEP_S (SHA512_F0o, SHA512_F1o, f, g, h, a, b, c, d, e, w3_t, k_sha512[i + 3]); \ + SHA512_STEP_S (SHA512_F0o, SHA512_F1o, e, f, g, h, a, b, c, d, w4_t, k_sha512[i + 4]); \ + SHA512_STEP_S (SHA512_F0o, SHA512_F1o, d, e, f, g, h, a, b, c, w5_t, k_sha512[i + 5]); \ + SHA512_STEP_S (SHA512_F0o, SHA512_F1o, c, d, e, f, g, h, a, b, w6_t, k_sha512[i + 6]); \ + SHA512_STEP_S (SHA512_F0o, SHA512_F1o, b, c, d, e, f, g, h, a, w7_t, k_sha512[i + 7]); \ + SHA512_STEP_S (SHA512_F0o, SHA512_F1o, a, b, c, d, e, f, g, h, w8_t, k_sha512[i + 8]); \ + SHA512_STEP_S (SHA512_F0o, SHA512_F1o, h, a, b, c, d, e, f, g, w9_t, k_sha512[i + 9]); \ + SHA512_STEP_S (SHA512_F0o, SHA512_F1o, g, h, a, b, c, d, e, f, wa_t, k_sha512[i + 10]); \ + SHA512_STEP_S (SHA512_F0o, SHA512_F1o, f, g, h, a, b, c, d, e, wb_t, k_sha512[i + 11]); \ + SHA512_STEP_S (SHA512_F0o, SHA512_F1o, e, f, g, h, a, b, c, d, wc_t, k_sha512[i + 12]); \ + SHA512_STEP_S (SHA512_F0o, SHA512_F1o, d, e, f, g, h, a, b, c, wd_t, k_sha512[i + 13]); \ + SHA512_STEP_S (SHA512_F0o, SHA512_F1o, c, d, e, f, g, h, a, b, we_t, k_sha512[i + 14]); \ + SHA512_STEP_S (SHA512_F0o, SHA512_F1o, b, c, d, e, f, g, h, a, wf_t, k_sha512[i + 15]); \ + } + + ROUND_STEP_S (0); + + #ifdef _unroll + #pragma unroll + #endif + for (int i = 16; i < 80; i += 16) + { + ROUND_EXPAND_S (); ROUND_STEP_S (i); + } + + digest[0] += a; + digest[1] += b; + digest[2] += c; + digest[3] += d; + digest[4] += e; + digest[5] += f; + digest[6] += g; + digest[7] += h; +} + +void hmac_sha512_pad_S (u64 w0[4], u64 w1[4], u64 w2[4], u64 w3[4], u64 ipad[8], u64 opad[8]) +{ + w0[0] = w0[0] ^ 0x3636363636363636; + w0[1] = w0[1] ^ 0x3636363636363636; + w0[2] = w0[2] ^ 0x3636363636363636; + w0[3] = w0[3] ^ 0x3636363636363636; + w1[0] = w1[0] ^ 0x3636363636363636; + w1[1] = w1[1] ^ 0x3636363636363636; + w1[2] = w1[2] ^ 0x3636363636363636; + w1[3] = w1[3] ^ 0x3636363636363636; + w2[0] = w2[0] ^ 0x3636363636363636; + w2[1] = w2[1] ^ 0x3636363636363636; + w2[2] = w2[2] ^ 0x3636363636363636; + w2[3] = w2[3] ^ 0x3636363636363636; + w3[0] = w3[0] ^ 0x3636363636363636; + w3[1] = w3[1] ^ 0x3636363636363636; + w3[2] = w3[2] ^ 0x3636363636363636; + w3[3] = w3[3] ^ 0x3636363636363636; + + ipad[0] = SHA512M_A; + ipad[1] = SHA512M_B; + ipad[2] = SHA512M_C; + ipad[3] = SHA512M_D; + ipad[4] = SHA512M_E; + ipad[5] = SHA512M_F; + ipad[6] = SHA512M_G; + ipad[7] = SHA512M_H; + + sha512_transform_S (w0, w1, w2, w3, ipad); + + w0[0] = w0[0] ^ 0x6a6a6a6a6a6a6a6a; + w0[1] = w0[1] ^ 0x6a6a6a6a6a6a6a6a; + w0[2] = w0[2] ^ 0x6a6a6a6a6a6a6a6a; + w0[3] = w0[3] ^ 0x6a6a6a6a6a6a6a6a; + w1[0] = w1[0] ^ 0x6a6a6a6a6a6a6a6a; + w1[1] = w1[1] ^ 0x6a6a6a6a6a6a6a6a; + w1[2] = w1[2] ^ 0x6a6a6a6a6a6a6a6a; + w1[3] = w1[3] ^ 0x6a6a6a6a6a6a6a6a; + w2[0] = w2[0] ^ 0x6a6a6a6a6a6a6a6a; + w2[1] = w2[1] ^ 0x6a6a6a6a6a6a6a6a; + w2[2] = w2[2] ^ 0x6a6a6a6a6a6a6a6a; + w2[3] = w2[3] ^ 0x6a6a6a6a6a6a6a6a; + w3[0] = w3[0] ^ 0x6a6a6a6a6a6a6a6a; + w3[1] = w3[1] ^ 0x6a6a6a6a6a6a6a6a; + w3[2] = w3[2] ^ 0x6a6a6a6a6a6a6a6a; + w3[3] = w3[3] ^ 0x6a6a6a6a6a6a6a6a; + + opad[0] = SHA512M_A; + opad[1] = SHA512M_B; + opad[2] = SHA512M_C; + opad[3] = SHA512M_D; + opad[4] = SHA512M_E; + opad[5] = SHA512M_F; + opad[6] = SHA512M_G; + opad[7] = SHA512M_H; + + sha512_transform_S (w0, w1, w2, w3, opad); +} + +void hmac_sha512_run_S (u64 w0[4], u64 w1[4], u64 w2[4], u64 w3[4], u64 ipad[8], u64 opad[8], u64 digest[8]) +{ + digest[0] = ipad[0]; + digest[1] = ipad[1]; + digest[2] = ipad[2]; + digest[3] = ipad[3]; + digest[4] = ipad[4]; + digest[5] = ipad[5]; + digest[6] = ipad[6]; + digest[7] = ipad[7]; + + sha512_transform_S (w0, w1, w2, w3, digest); + + w0[0] = digest[0]; + w0[1] = digest[1]; + w0[2] = digest[2]; + w0[3] = digest[3]; + w1[0] = digest[4]; + w1[1] = digest[5]; + w1[2] = digest[6]; + w1[3] = digest[7]; + w2[0] = 0x8000000000000000; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = (128 + 64) * 8; + + digest[0] = opad[0]; + digest[1] = opad[1]; + digest[2] = opad[2]; + digest[3] = opad[3]; + digest[4] = opad[4]; + digest[5] = opad[5]; + digest[6] = opad[6]; + digest[7] = opad[7]; + + sha512_transform_S (w0, w1, w2, w3, digest); +} + +void sha512_transform_V (const u64x w0[4], const u64x w1[4], const u64x w2[4], const u64x w3[4], u64x digest[8]) +{ + u64x a = digest[0]; + u64x b = digest[1]; + u64x c = digest[2]; + u64x d = digest[3]; + u64x e = digest[4]; + u64x f = digest[5]; + u64x g = digest[6]; + u64x h = digest[7]; + + u64x w0_t = w0[0]; + u64x w1_t = w0[1]; + u64x w2_t = w0[2]; + u64x w3_t = w0[3]; + u64x w4_t = w1[0]; + u64x w5_t = w1[1]; + u64x w6_t = w1[2]; + u64x w7_t = w1[3]; + u64x w8_t = w2[0]; + u64x w9_t = w2[1]; + u64x wa_t = w2[2]; + u64x wb_t = w2[3]; + u64x wc_t = w3[0]; + u64x wd_t = w3[1]; + u64x we_t = w3[2]; + u64x wf_t = w3[3]; + + #define ROUND_EXPAND() \ + { \ + w0_t = SHA512_EXPAND (we_t, w9_t, w1_t, w0_t); \ + w1_t = SHA512_EXPAND (wf_t, wa_t, w2_t, w1_t); \ + w2_t = SHA512_EXPAND (w0_t, wb_t, w3_t, w2_t); \ + w3_t = SHA512_EXPAND (w1_t, wc_t, w4_t, w3_t); \ + w4_t = SHA512_EXPAND (w2_t, wd_t, w5_t, w4_t); \ + w5_t = SHA512_EXPAND (w3_t, we_t, w6_t, w5_t); \ + w6_t = SHA512_EXPAND (w4_t, wf_t, w7_t, w6_t); \ + w7_t = SHA512_EXPAND (w5_t, w0_t, w8_t, w7_t); \ + w8_t = SHA512_EXPAND (w6_t, w1_t, w9_t, w8_t); \ + w9_t = SHA512_EXPAND (w7_t, w2_t, wa_t, w9_t); \ + wa_t = SHA512_EXPAND (w8_t, w3_t, wb_t, wa_t); \ + wb_t = SHA512_EXPAND (w9_t, w4_t, wc_t, wb_t); \ + wc_t = SHA512_EXPAND (wa_t, w5_t, wd_t, wc_t); \ + wd_t = SHA512_EXPAND (wb_t, w6_t, we_t, wd_t); \ + we_t = SHA512_EXPAND (wc_t, w7_t, wf_t, we_t); \ + wf_t = SHA512_EXPAND (wd_t, w8_t, w0_t, wf_t); \ + } + + #define ROUND_STEP(i) \ + { \ + SHA512_STEP (SHA512_F0o, SHA512_F1o, a, b, c, d, e, f, g, h, w0_t, k_sha512[i + 0]); \ + SHA512_STEP (SHA512_F0o, SHA512_F1o, h, a, b, c, d, e, f, g, w1_t, k_sha512[i + 1]); \ + SHA512_STEP (SHA512_F0o, SHA512_F1o, g, h, a, b, c, d, e, f, w2_t, k_sha512[i + 2]); \ + SHA512_STEP (SHA512_F0o, SHA512_F1o, f, g, h, a, b, c, d, e, w3_t, k_sha512[i + 3]); \ + SHA512_STEP (SHA512_F0o, SHA512_F1o, e, f, g, h, a, b, c, d, w4_t, k_sha512[i + 4]); \ + SHA512_STEP (SHA512_F0o, SHA512_F1o, d, e, f, g, h, a, b, c, w5_t, k_sha512[i + 5]); \ + SHA512_STEP (SHA512_F0o, SHA512_F1o, c, d, e, f, g, h, a, b, w6_t, k_sha512[i + 6]); \ + SHA512_STEP (SHA512_F0o, SHA512_F1o, b, c, d, e, f, g, h, a, w7_t, k_sha512[i + 7]); \ + SHA512_STEP (SHA512_F0o, SHA512_F1o, a, b, c, d, e, f, g, h, w8_t, k_sha512[i + 8]); \ + SHA512_STEP (SHA512_F0o, SHA512_F1o, h, a, b, c, d, e, f, g, w9_t, k_sha512[i + 9]); \ + SHA512_STEP (SHA512_F0o, SHA512_F1o, g, h, a, b, c, d, e, f, wa_t, k_sha512[i + 10]); \ + SHA512_STEP (SHA512_F0o, SHA512_F1o, f, g, h, a, b, c, d, e, wb_t, k_sha512[i + 11]); \ + SHA512_STEP (SHA512_F0o, SHA512_F1o, e, f, g, h, a, b, c, d, wc_t, k_sha512[i + 12]); \ + SHA512_STEP (SHA512_F0o, SHA512_F1o, d, e, f, g, h, a, b, c, wd_t, k_sha512[i + 13]); \ + SHA512_STEP (SHA512_F0o, SHA512_F1o, c, d, e, f, g, h, a, b, we_t, k_sha512[i + 14]); \ + SHA512_STEP (SHA512_F0o, SHA512_F1o, b, c, d, e, f, g, h, a, wf_t, k_sha512[i + 15]); \ + } + + ROUND_STEP (0); + + #ifdef _unroll + #pragma unroll + #endif + for (int i = 16; i < 80; i += 16) + { + ROUND_EXPAND (); ROUND_STEP (i); + } + + digest[0] += a; + digest[1] += b; + digest[2] += c; + digest[3] += d; + digest[4] += e; + digest[5] += f; + digest[6] += g; + digest[7] += h; +} + +void hmac_sha512_run_V (u64x w0[4], u64x w1[4], u64x w2[4], u64x w3[4], u64x ipad[8], u64x opad[8], u64x digest[8]) +{ + digest[0] = ipad[0]; + digest[1] = ipad[1]; + digest[2] = ipad[2]; + digest[3] = ipad[3]; + digest[4] = ipad[4]; + digest[5] = ipad[5]; + digest[6] = ipad[6]; + digest[7] = ipad[7]; + + sha512_transform_V (w0, w1, w2, w3, digest); + + w0[0] = digest[0]; + w0[1] = digest[1]; + w0[2] = digest[2]; + w0[3] = digest[3]; + w1[0] = digest[4]; + w1[1] = digest[5]; + w1[2] = digest[6]; + w1[3] = digest[7]; + w2[0] = 0x8000000000000000; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = (128 + 64) * 8; + + digest[0] = opad[0]; + digest[1] = opad[1]; + digest[2] = opad[2]; + digest[3] = opad[3]; + digest[4] = opad[4]; + digest[5] = opad[5]; + digest[6] = opad[6]; + digest[7] = opad[7]; + + sha512_transform_V (w0, w1, w2, w3, digest); +} + +__kernel void m15300_init (__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) +{ + /** + * base + */ + + const u32 gid = get_global_id (0); + + if (gid >= gid_max) return; + + u32 w0[4]; + + w0[0] = pws[gid].i[0]; + w0[1] = pws[gid].i[1]; + w0[2] = pws[gid].i[2]; + w0[3] = pws[gid].i[3]; + + u32 w1[4]; + + w1[0] = pws[gid].i[4]; + w1[1] = pws[gid].i[5]; + w1[2] = pws[gid].i[6]; + w1[3] = pws[gid].i[7]; + + u32 w2[4]; + + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + + u32 w3[4]; + + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + const u32 pw_len = pws[gid].pw_len; + + append_0x80_2x4_S (w0, w1, pw_len); + + make_unicode_S (w1, w2, w3); + make_unicode_S (w0, w0, w1); + + /** + * main + */ + + /** + * salt == SID + */ + + const u32 salt_len = esalt_bufs[digests_offset].SID_len; + + u32 digest_context[5]; + + /* local credentials */ + if (esalt_bufs[digests_offset].context == 1) + { + digest_context[0] = SHA1M_A; + digest_context[1] = SHA1M_B; + digest_context[2] = SHA1M_C; + digest_context[3] = SHA1M_D; + digest_context[4] = SHA1M_E; + + w0[0] = swap32_S (w0[0]); + w0[1] = swap32_S (w0[1]); + w0[2] = swap32_S (w0[2]); + w0[3] = swap32_S (w0[3]); + + w1[0] = swap32_S (w1[0]); + w1[1] = swap32_S (w1[1]); + w1[2] = swap32_S (w1[2]); + w1[3] = swap32_S (w1[3]); + + w2[0] = swap32_S (w2[0]); + w2[1] = swap32_S (w2[1]); + w2[2] = swap32_S (w2[2]); + w2[3] = swap32_S (w2[3]); + + w3[0] = swap32_S (w3[0]); + w3[1] = swap32_S (w3[1]); + w3[2] = 0; + w3[3] = pw_len * 2 * 8; + + sha1_transform_S (w0, w1, w2, w3, digest_context); + } + /* domain credentials */ + else if (esalt_bufs[digests_offset].context == 2) + { + digest_context[0] = MD4M_A; + digest_context[1] = MD4M_B; + digest_context[2] = MD4M_C; + digest_context[3] = MD4M_D; + + w3[2] = pw_len * 2 * 8; + + md4_transform_S (w0, w1, w2, w3, digest_context); + + digest_context[0] = swap32_S (digest_context[0]); + digest_context[1] = swap32_S (digest_context[1]); + digest_context[2] = swap32_S (digest_context[2]); + digest_context[3] = swap32_S (digest_context[3]); + digest_context[4] = 0; + } + /* initialize hmac-sha1 */ + + /** + * pads + */ + + w0[0] = digest_context[0]; + w0[1] = digest_context[1]; + w0[2] = digest_context[2]; + w0[3] = digest_context[3]; + w1[0] = digest_context[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; + + u32 ipad[5]; + u32 opad[5]; + + hmac_sha1_pad_S (w0, w1, w2, w3, ipad, opad); + + tmps[gid].ipad[0] = ipad[0]; + tmps[gid].ipad[1] = ipad[1]; + tmps[gid].ipad[2] = ipad[2]; + tmps[gid].ipad[3] = ipad[3]; + tmps[gid].ipad[4] = ipad[4]; + + tmps[gid].opad[0] = opad[0]; + tmps[gid].opad[1] = opad[1]; + tmps[gid].opad[2] = opad[2]; + tmps[gid].opad[3] = opad[3]; + tmps[gid].opad[4] = opad[4]; + + /** + * hmac1 + */ + + w0[0] = esalt_bufs[digests_offset].SID[ 0]; + w0[1] = esalt_bufs[digests_offset].SID[ 1]; + w0[2] = esalt_bufs[digests_offset].SID[ 2]; + w0[3] = esalt_bufs[digests_offset].SID[ 3]; + w1[0] = esalt_bufs[digests_offset].SID[ 4]; + w1[1] = esalt_bufs[digests_offset].SID[ 5]; + w1[2] = esalt_bufs[digests_offset].SID[ 6]; + w1[3] = esalt_bufs[digests_offset].SID[ 7]; + w2[0] = esalt_bufs[digests_offset].SID[ 8]; + w2[1] = esalt_bufs[digests_offset].SID[ 9]; + w2[2] = esalt_bufs[digests_offset].SID[10]; + w2[3] = esalt_bufs[digests_offset].SID[11]; + w3[0] = esalt_bufs[digests_offset].SID[12]; + w3[1] = esalt_bufs[digests_offset].SID[13]; + w3[2] = esalt_bufs[digests_offset].SID[14]; + w3[3] = esalt_bufs[digests_offset].SID[15]; + + sha1_transform_S (w0, w1, w2, w3, ipad); + + w0[0] = esalt_bufs[digests_offset].SID[16 + 0]; + w0[1] = esalt_bufs[digests_offset].SID[16 + 1]; + w0[2] = esalt_bufs[digests_offset].SID[16 + 2]; + w0[3] = esalt_bufs[digests_offset].SID[16 + 3]; + w1[0] = esalt_bufs[digests_offset].SID[16 + 4]; + w1[1] = esalt_bufs[digests_offset].SID[16 + 5]; + w1[2] = esalt_bufs[digests_offset].SID[16 + 6]; + w1[3] = esalt_bufs[digests_offset].SID[16 + 7]; + w2[0] = esalt_bufs[digests_offset].SID[16 + 8]; + w2[1] = esalt_bufs[digests_offset].SID[16 + 9]; + w2[2] = esalt_bufs[digests_offset].SID[16 + 10]; + w2[3] = esalt_bufs[digests_offset].SID[16 + 11]; + w3[0] = esalt_bufs[digests_offset].SID[16 + 12]; + w3[1] = esalt_bufs[digests_offset].SID[16 + 13]; + w3[2] = 0; + w3[3] = (64 + salt_len) * 8; + + u32 key[5]; + + hmac_sha1_run_S (w0, w1, w2, w3, ipad, opad, key); + /* this key is used as password for pbkdf2-hmac-* */ + + tmps[gid].userKey[0] = key[0]; + tmps[gid].userKey[1] = key[1]; + tmps[gid].userKey[2] = key[2]; + tmps[gid].userKey[3] = key[3]; + tmps[gid].userKey[4] = key[4]; + + /* if DPAPImk version 1, pbkdf-hmac-sha1 is used */ + if (esalt_bufs[digests_offset].version == 1) + { + /** + * pads + */ + + w0[0] = key[0]; + w0[1] = key[1]; + w0[2] = key[2]; + w0[3] = key[3]; + w1[0] = key[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); + + tmps[gid].ipad[0] = ipad[0]; + tmps[gid].ipad[1] = ipad[1]; + tmps[gid].ipad[2] = ipad[2]; + tmps[gid].ipad[3] = ipad[3]; + tmps[gid].ipad[4] = ipad[4]; + + tmps[gid].opad[0] = opad[0]; + tmps[gid].opad[1] = opad[1]; + tmps[gid].opad[2] = opad[2]; + tmps[gid].opad[3] = opad[3]; + tmps[gid].opad[4] = opad[4]; + + /** + * hmac1 + */ + + 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; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = (64 + 16 + 4) * 8; + + u32 digest[5]; + + hmac_sha1_run_S (w0, w1, w2, w3, ipad, opad, digest); + + tmps[gid].dgst[0] = digest[0]; + tmps[gid].dgst[1] = digest[1]; + tmps[gid].dgst[2] = digest[2]; + tmps[gid].dgst[3] = digest[3]; + tmps[gid].dgst[4] = digest[4]; + + tmps[gid].out[0] = digest[0]; + tmps[gid].out[1] = digest[1]; + 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) + { + u64 w0_x64[4]; + u64 w1_x64[4]; + u64 w2_x64[4]; + u64 w3_x64[4]; + + w0_x64[0] = hl32_to_64_S (key[0], key[1]); + w0_x64[1] = hl32_to_64_S (key[2], key[3]); + w0_x64[2] = hl32_to_64_S (key[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); + + tmps[gid].ipad64[0] = ipad64[0]; + tmps[gid].ipad64[1] = ipad64[1]; + tmps[gid].ipad64[2] = ipad64[2]; + tmps[gid].ipad64[3] = ipad64[3]; + tmps[gid].ipad64[4] = ipad64[4]; + tmps[gid].ipad64[5] = ipad64[5]; + tmps[gid].ipad64[6] = ipad64[6]; + tmps[gid].ipad64[7] = ipad64[7]; + + tmps[gid].opad64[0] = opad64[0]; + tmps[gid].opad64[1] = opad64[1]; + tmps[gid].opad64[2] = opad64[2]; + tmps[gid].opad64[3] = opad64[3]; + tmps[gid].opad64[4] = opad64[4]; + tmps[gid].opad64[5] = opad64[5]; + tmps[gid].opad64[6] = opad64[6]; + tmps[gid].opad64[7] = opad64[7]; + + 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; + 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 + 4) * 8; + + u64 dgst64[8]; + + hmac_sha512_run_S (w0_x64, w1_x64, w2_x64, w3_x64, ipad64, opad64, dgst64); + + tmps[gid].dgst64[0] = dgst64[0]; + tmps[gid].dgst64[1] = dgst64[1]; + tmps[gid].dgst64[2] = dgst64[2]; + tmps[gid].dgst64[3] = dgst64[3]; + tmps[gid].dgst64[4] = dgst64[4]; + tmps[gid].dgst64[5] = dgst64[5]; + tmps[gid].dgst64[6] = dgst64[6]; + tmps[gid].dgst64[7] = dgst64[7]; + + tmps[gid].out64[0] = dgst64[0]; + tmps[gid].out64[1] = dgst64[1]; + tmps[gid].out64[2] = dgst64[2]; + tmps[gid].out64[3] = dgst64[3]; + tmps[gid].out64[4] = dgst64[4]; + tmps[gid].out64[5] = dgst64[5]; + tmps[gid].out64[6] = dgst64[6]; + tmps[gid].out64[7] = dgst64[7]; + } +} + +__kernel void m15300_loop (__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) +{ + /** + * base + */ + + const u32 gid = get_global_id (0); + + if ((gid * VECT_SIZE) >= gid_max) return; + + if (esalt_bufs[digests_offset].version == 1) + { + u32x ipad[5]; + u32x opad[5]; + + ipad[0] = packv (tmps, ipad, gid, 0); + ipad[1] = packv (tmps, ipad, gid, 1); + ipad[2] = packv (tmps, ipad, gid, 2); + ipad[3] = packv (tmps, ipad, gid, 3); + ipad[4] = packv (tmps, ipad, gid, 4); + + opad[0] = packv (tmps, opad, gid, 0); + opad[1] = packv (tmps, opad, gid, 1); + opad[2] = packv (tmps, opad, gid, 2); + opad[3] = packv (tmps, opad, gid, 3); + opad[4] = packv (tmps, opad, gid, 4); + + /** + * iter1 + */ + + for (u32 i = 0; i < 8; i += 5) + { + 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); + + 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]); + } + } + else if (esalt_bufs[digests_offset].version == 2) + { + u64x ipad[8]; + u64x opad[8]; + + ipad[0] = pack64v (tmps, ipad64, gid, 0); + ipad[1] = pack64v (tmps, ipad64, gid, 1); + ipad[2] = pack64v (tmps, ipad64, gid, 2); + ipad[3] = pack64v (tmps, ipad64, gid, 3); + ipad[4] = pack64v (tmps, ipad64, gid, 4); + ipad[5] = pack64v (tmps, ipad64, gid, 5); + ipad[6] = pack64v (tmps, ipad64, gid, 6); + ipad[7] = pack64v (tmps, ipad64, gid, 7); + + opad[0] = pack64v (tmps, opad64, gid, 0); + opad[1] = pack64v (tmps, opad64, gid, 1); + opad[2] = pack64v (tmps, opad64, gid, 2); + opad[3] = pack64v (tmps, opad64, gid, 3); + opad[4] = pack64v (tmps, opad64, gid, 4); + opad[5] = pack64v (tmps, opad64, gid, 5); + opad[6] = pack64v (tmps, opad64, gid, 6); + opad[7] = pack64v (tmps, opad64, gid, 7); + + u64x dgst[8]; + u64x out[8]; + + dgst[0] = pack64v (tmps, dgst64, gid, 0); + dgst[1] = pack64v (tmps, dgst64, gid, 1); + dgst[2] = pack64v (tmps, dgst64, gid, 2); + dgst[3] = pack64v (tmps, dgst64, gid, 3); + dgst[4] = pack64v (tmps, dgst64, gid, 4); + dgst[5] = pack64v (tmps, dgst64, gid, 5); + dgst[6] = pack64v (tmps, dgst64, gid, 6); + dgst[7] = pack64v (tmps, dgst64, gid, 7); + + out[0] = pack64v (tmps, out64, gid, 0); + out[1] = pack64v (tmps, out64, gid, 1); + out[2] = pack64v (tmps, out64, gid, 2); + out[3] = pack64v (tmps, out64, gid, 3); + out[4] = pack64v (tmps, out64, gid, 4); + out[5] = pack64v (tmps, out64, gid, 5); + out[6] = pack64v (tmps, out64, gid, 6); + out[7] = pack64v (tmps, out64, gid, 7); + + for (u32 j = 0; j < loop_cnt; j++) + { + u64x w0[4]; + u64x w1[4]; + u64x w2[4]; + u64x 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] = out[5]; + w1[2] = out[6]; + w1[3] = out[7]; + w2[0] = 0x8000000000000000; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = (128 + 64) * 8; + + hmac_sha512_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]; + out[5] ^= dgst[5]; + out[6] ^= dgst[6]; + out[7] ^= dgst[7]; + } + + unpackv (tmps, dgst64, gid, 0, dgst[0]); + unpackv (tmps, dgst64, gid, 1, dgst[1]); + unpackv (tmps, dgst64, gid, 2, dgst[2]); + unpackv (tmps, dgst64, gid, 3, dgst[3]); + unpackv (tmps, dgst64, gid, 4, dgst[4]); + unpackv (tmps, dgst64, gid, 5, dgst[5]); + unpackv (tmps, dgst64, gid, 6, dgst[6]); + unpackv (tmps, dgst64, gid, 7, dgst[7]); + + unpackv (tmps, out64, gid, 0, out[0]); + unpackv (tmps, out64, gid, 1, out[1]); + unpackv (tmps, out64, gid, 2, out[2]); + unpackv (tmps, out64, gid, 3, out[3]); + unpackv (tmps, out64, gid, 4, out[4]); + unpackv (tmps, out64, gid, 5, out[5]); + unpackv (tmps, out64, gid, 6, out[6]); + unpackv (tmps, out64, gid, 7, out[7]); + } +} + +__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) +{ + 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]); + + u32 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 u32 a = (key[0]); + const u32 b = (key[1]); + + u32 Ka[16]; + u32 Kb[16]; + + _des_crypt_keysetup (a, b, Ka, Kb, s_skb); + + const u32 c = (key[2]); + const u32 d = (key[3]); + + u32 Kc[16]; + u32 Kd[16]; + + _des_crypt_keysetup (c, d, Kc, Kd, s_skb); + + const u32 e = (key[4]); + const u32 f = (key[5]); + + u32 Ke[16]; + u32 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 */ + + u32 data[2]; + + data[0] = swap32_S (esalt_bufs[digests_offset].contents[contents_off + 0]); + data[1] = swap32_S (esalt_bufs[digests_offset].contents[contents_off + 1]); + + u32 p1[2]; + + _des_crypt_decrypt (p1, data, Ke, Kf, s_SPtrans); + + /* Second Pass */ + + u32 p2[2]; + + _des_crypt_encrypt (p2, p1, Kc, Kd, s_SPtrans); + + /* Third Pass */ + + u32 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 expectedHmac[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]); + + expectedHmac[0] = swap32_S (decrypted[4 + 0]); + expectedHmac[1] = swap32_S (decrypted[4 + 1]); + expectedHmac[2] = swap32_S (decrypted[4 + 2]); + expectedHmac[3] = swap32_S (decrypted[4 + 3]); + expectedHmac[4] = swap32_S (decrypted[4 + 4]); + + for(int i = 0; i < 16; i++) + { + lastKey[i] = decrypted[i + 26 - 16]; + } + + w0[0] = tmps[gid].userKey[0]; + w0[1] = tmps[gid].userKey[1]; + w0[2] = tmps[gid].userKey[2]; + w0[3] = tmps[gid].userKey[3]; + w1[0] = tmps[gid].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 ( expectedHmac[0] == digest[0] + && expectedHmac[1] == digest[1] + && expectedHmac[2] == digest[2] + && expectedHmac[3] == digest[3] + && expectedHmac[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_S (tmps[gid].out64[0]); + key[1] = l32_from_64_S (tmps[gid].out64[0]); + key[2] = h32_from_64_S (tmps[gid].out64[1]); + key[3] = l32_from_64_S (tmps[gid].out64[1]); + key[4] = h32_from_64_S (tmps[gid].out64[2]); + key[5] = l32_from_64_S (tmps[gid].out64[2]); + key[6] = h32_from_64_S (tmps[gid].out64[3]); + key[7] = l32_from_64_S (tmps[gid].out64[3]); + + u32 iv[4]; + + iv[0] = h32_from_64_S (tmps[gid].out64[4]); + iv[1] = l32_from_64_S (tmps[gid].out64[4]); + iv[2] = h32_from_64_S (tmps[gid].out64[5]); + iv[3] = l32_from_64_S (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 expectedHmac[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++) + { + expectedHmac[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 (tmps[gid].userKey[0], tmps[gid].userKey[1]); + w0_x64[1] = hl32_to_64_S (tmps[gid].userKey[2], tmps[gid].userKey[3]); + w0_x64[2] = hl32_to_64_S (tmps[gid].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 ( expectedHmac[ 0] == h32_from_64_S (dgst64[0]) + && expectedHmac[ 1] == l32_from_64_S (dgst64[0]) + && expectedHmac[ 2] == h32_from_64_S (dgst64[1]) + && expectedHmac[ 3] == l32_from_64_S (dgst64[1]) + && expectedHmac[12] == h32_from_64_S (dgst64[6]) + && expectedHmac[13] == l32_from_64_S (dgst64[6]) + && expectedHmac[14] == h32_from_64_S (dgst64[7]) + && expectedHmac[15] == l32_from_64_S (dgst64[7])) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos); + } + } +} diff --git a/docs/changes.txt b/docs/changes.txt index 7acfae9e8..3757de308 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -6,6 +6,7 @@ - Added hash-mode 600 = Blake2-512 - Added hash-mode 15200 = Blockchain, My Wallet, V2 +- Added hash-mode 15300 = DPAPI masterkey file v1 and v2 ## ## Features diff --git a/docs/credits.txt b/docs/credits.txt index dbb1b6c51..db0a6855d 100644 --- a/docs/credits.txt +++ b/docs/credits.txt @@ -31,6 +31,7 @@ Jean-Christophe "Fist0urs" Delaunay (@F * Kerberos TGS Rep enctype 23 kernel module * AxCrypt kernel module * KeePass kernel module +* DPAPImk v1 and v2 kernel module Other contributors to hashcat diff --git a/docs/readme.txt b/docs/readme.txt index d128e38e6..22007843a 100644 --- a/docs/readme.txt +++ b/docs/readme.txt @@ -145,6 +145,7 @@ NVIDIA GPUs require "NVIDIA Driver" (367.x or later) - NTLM - Domain Cached Credentials (DCC), MS Cache - Domain Cached Credentials 2 (DCC2), MS Cache 2 +- DPAPI masterkey file v1 and v2 - MS-AzureSync PBKDF2-HMAC-SHA256 - descrypt - bsdicrypt diff --git a/extra/tab_completion/hashcat.sh b/extra/tab_completion/hashcat.sh index 11d07c76a..261982863 100644 --- a/extra/tab_completion/hashcat.sh +++ b/extra/tab_completion/hashcat.sh @@ -176,7 +176,7 @@ _hashcat () { local VERSION=3.5.0 - local HASH_MODES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 124 130 131 132 133 140 141 150 160 200 300 400 500 501 600 900 1000 1100 1400 1410 1411 1420 1421 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 3910 4010 4110 4300 4400 4500 4520 4521 4522 4700 4800 4900 5000 5100 5200 5300 5400 5500 5600 5700 5800 6000 6100 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6300 6400 6500 6600 6700 6800 6900 7000 7100 7200 7300 7400 7500 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8800 8900 9000 9100 9200 9300 9400 9500 9600 9700 9710 9720 9800 9810 9820 9900 10000 10100 10200 10300 10400 10410 10420 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11700 11800 11900 12000 12001 12100 12200 12300 12400 12500 12600 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 13900 14000 14100 14700 14800 14900 15000 15100 15200" + local HASH_MODES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 124 130 131 132 133 140 141 150 160 200 300 400 500 501 600 900 1000 1100 1400 1410 1411 1420 1421 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 3910 4010 4110 4300 4400 4500 4520 4521 4522 4700 4800 4900 5000 5100 5200 5300 5400 5500 5600 5700 5800 6000 6100 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6300 6400 6500 6600 6700 6800 6900 7000 7100 7200 7300 7400 7500 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8800 8900 9000 9100 9200 9300 9400 9500 9600 9700 9710 9720 9800 9810 9820 9900 10000 10100 10200 10300 10400 10410 10420 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11700 11800 11900 12000 12001 12100 12200 12300 12400 12500 12600 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 13900 14000 14100 14700 14800 14900 15000 15100 15200 15300" local ATTACK_MODES="0 1 3 6 7" local HCCAPX_MESSAGE_PAIR="0 1 2 3 4 5" local OUTFILE_FORMATS="1 2 3 4 5 6 7 8 9 10 11 12 13 14 15" diff --git a/include/interface.h b/include/interface.h index e9c6fbc2a..4765b07e2 100644 --- a/include/interface.h +++ b/include/interface.h @@ -49,28 +49,28 @@ static const char LM_MASKED_PLAIN[] = "[notfound]"; #define LUKS_ALIGN_KEYSLOTS 4096 struct luks_phdr { - char magic[LUKS_MAGIC_L]; - uint16_t version; - char cipherName[LUKS_CIPHERNAME_L]; - char cipherMode[LUKS_CIPHERMODE_L]; - char hashSpec[LUKS_HASHSPEC_L]; - uint32_t payloadOffset; - uint32_t keyBytes; - char mkDigest[LUKS_DIGESTSIZE]; - char mkDigestSalt[LUKS_SALTSIZE]; - uint32_t mkDigestIterations; - char uuid[UUID_STRING_L]; - struct { - uint32_t active; - /* parameters used for password processing */ - uint32_t passwordIterations; - char passwordSalt[LUKS_SALTSIZE]; - /* parameters used for AF store/load */ - uint32_t keyMaterialOffset; - uint32_t stripes; - } keyblock[LUKS_NUMKEYS]; - /* Align it to 512 sector size */ - char _padding[432]; + char magic[LUKS_MAGIC_L]; + uint16_t version; + char cipherName[LUKS_CIPHERNAME_L]; + char cipherMode[LUKS_CIPHERMODE_L]; + char hashSpec[LUKS_HASHSPEC_L]; + uint32_t payloadOffset; + uint32_t keyBytes; + char mkDigest[LUKS_DIGESTSIZE]; + char mkDigestSalt[LUKS_SALTSIZE]; + uint32_t mkDigestIterations; + char uuid[UUID_STRING_L]; + struct { + uint32_t active; + /* parameters used for password processing */ + uint32_t passwordIterations; + char passwordSalt[LUKS_SALTSIZE]; + /* parameters used for AF store/load */ + uint32_t keyMaterialOffset; + uint32_t stripes; + } keyblock[LUKS_NUMKEYS]; + /* Align it to 512 sector size */ + char _padding[432]; }; // not from original headers start with hc_ @@ -420,6 +420,27 @@ typedef struct psafe3 } psafe3_t; +typedef struct dpapimk +{ + u32 version; + u32 context; + + u32 SID[32]; + u32 SID_len; + u32 SID_offset; + + /* here only for possible + forward compatibiliy + */ + // u8 cipher_algo[16]; + // u8 hash_algo[16]; + + u32 iv[4]; + u32 contents_len; + u32 contents[128]; + +} dpapimk_t; + typedef struct pdf14_tmp { u32 digest[4]; @@ -810,6 +831,24 @@ typedef struct keepass_tmp } keepass_tmp_t; +typedef struct dpapimk_tmp +{ + /* dedicated to hmac-sha1 */ + u32 ipad[5]; + u32 opad[5]; + u32 dgst[10]; + u32 out[10]; + + u32 userKey[5]; + + /* dedicated to hmac-sha512 */ + u64 ipad64[8]; + u64 opad64[8]; + u64 dgst64[16]; + u64 out64[16]; + +} dpapimk_tmp_t; + typedef struct seven_zip_hook { u32 ukey[8]; @@ -1190,6 +1229,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, + DISPLAY_LEN_MIN_15300 = 1 + 7 + 1 + 1 + 1 + 1 + 1 + 10 + 1 + 4 + 1 + 4 + 1 + 1 + 1 + 32 + 1 + 3 + 1 + 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, DISPLAY_LEN_MIN_99999 = 1, DISPLAY_LEN_MAX_99999 = 55, @@ -1329,6 +1370,8 @@ typedef enum hash_type HASH_TYPE_ITUNES_BACKUP_10 = 57, HASH_TYPE_SKIP32 = 58, HASH_TYPE_BLAKE2B = 59, + HASH_TYPE_CHACHA20 = 60, + HASH_TYPE_DPAPIMK = 61, } hash_type_t; @@ -1518,6 +1561,7 @@ typedef enum kern_type KERN_TYPE_SKIP32 = 14900, KERN_TYPE_FILEZILLA_SERVER = 15000, KERN_TYPE_NETBSD_SHA1CRYPT = 15100, + KERN_TYPE_DPAPIMK = 15300, KERN_TYPE_PLAINTEXT = 99999, } kern_type_t; @@ -1588,6 +1632,7 @@ typedef enum rounds_count ROUNDS_ITUNES102_BACKUP = 10000, ROUNDS_ATLASSIAN = 10000, ROUNDS_NETBSD_SHA1CRYPT = 20000, + ROUNDS_DPAPIMK = 24000 - 1, // from 4000 to 24000 (possibly more) ROUNDS_STDOUT = 0 } rounds_count_t; @@ -1766,6 +1811,7 @@ 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); +int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig); /** * hook functions diff --git a/src/benchmark.c b/src/benchmark.c index fd85e023b..32774d016 100644 --- a/src/benchmark.c +++ b/src/benchmark.c @@ -6,7 +6,7 @@ #include "common.h" #include "benchmark.h" -const unsigned int DEFAULT_BENCHMARK_ALGORITHMS_CNT = 150; +const unsigned int DEFAULT_BENCHMARK_ALGORITHMS_CNT = 151; const unsigned int DEFAULT_BENCHMARK_ALGORITHMS_BUF[] = { @@ -85,6 +85,7 @@ const unsigned int DEFAULT_BENCHMARK_ALGORITHMS_BUF[] = 1000, 1100, 2100, + 15300, 12800, 1500, 12400, diff --git a/src/interface.c b/src/interface.c index 5b040df89..2060afb54 100644 --- a/src/interface.c +++ b/src/interface.c @@ -240,6 +240,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"; +static const char HT_15300[] = "DPAPI masterkey file v1 and v2"; static const char HT_99999[] = "Plaintext"; static const char HT_00011[] = "Joomla < 2.5.18"; @@ -317,6 +318,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$"; +static const char SIGNATURE_DPAPIMK[] = "$DPAPImk$"; static const char SIGNATURE_DRUPAL7[] = "$S$"; static const char SIGNATURE_ECRYPTFS[] = "$ecryptfs$"; static const char SIGNATURE_EPISERVER4[] = "$episerver$*1*"; @@ -2831,6 +2833,164 @@ int dcc2_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSE return (PARSER_OK); } +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); + + if (memcmp (SIGNATURE_DPAPIMK, input_buf, 9)) return (PARSER_SIGNATURE_UNMATCHED); + + u32 *digest = (u32 *) hash_buf->digest; + + salt_t *salt = hash_buf->salt; + + dpapimk_t *dpapimk = (dpapimk_t *) hash_buf->esalt; + + /** + * parse line + */ + + u8 *version_pos; + u8 *context_pos; + u8 *SID_pos; + u8 *cipher_algo_pos; // here just for possible forward compatibilities + u8 *hash_algo_pos; // same + u8 *rounds_pos; + u32 iv_len = 32; + u8 *iv_pos; + u8 *contents_len_pos; + u8 *contents_pos; + + version_pos = input_buf + 8 + 1; + + context_pos = (u8 *) strchr ((const char *) version_pos, '*'); + + if (context_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED); + + context_pos++; + + SID_pos = (u8 *) strchr ((const char *) context_pos, '*'); + + if (SID_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED); + + SID_pos++; + + cipher_algo_pos = (u8 *) strchr ((const char *) SID_pos, '*'); + + if (cipher_algo_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED); + + cipher_algo_pos++; + + hash_algo_pos = (u8 *) strchr ((const char *) cipher_algo_pos, '*'); + + if (hash_algo_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED); + + hash_algo_pos++; + + rounds_pos = (u8 *) strchr ((const char *) hash_algo_pos, '*'); + + if (rounds_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED); + + rounds_pos++; + + iv_pos = (u8 *) strchr ((const char *) rounds_pos, '*'); + + if (iv_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED); + + iv_pos++; + + contents_len_pos = (u8 *) strchr ((const char *) iv_pos, '*'); + + if (contents_len_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED); + + if (contents_len_pos - iv_pos != iv_len) return (PARSER_SALT_LENGTH); + + if (is_valid_hex_string (iv_pos, 32) == false) return (PARSER_SALT_ENCODING); + + contents_len_pos++; + + contents_pos = (u8 *) strchr ((const char *) contents_len_pos, '*'); + + if (contents_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED); + + contents_pos++; + + u32 version = atoll ((const char *) version_pos); + u32 contents_len = atoll ((const char *) contents_len_pos); + + if (version == 1 && contents_len != 208) return (PARSER_SALT_LENGTH); + if (version == 2 && contents_len != 288) return (PARSER_SALT_LENGTH); + + if (is_valid_hex_string (contents_pos, contents_len) == false) return (PARSER_SALT_ENCODING); + + u8 *end_line = (u8 *) strchr ((const char *) contents_pos, 0); + + if (end_line - contents_pos != contents_len) return (PARSER_SALT_LENGTH); + + dpapimk->version = version; + + dpapimk->context = atoll ((const char *) context_pos); + + salt->salt_iter = (atoll ((const char *) rounds_pos)) - 1; + + dpapimk->iv[0] = hex_to_u32 ((const u8 *) &iv_pos[ 0]); + dpapimk->iv[1] = hex_to_u32 ((const u8 *) &iv_pos[ 8]); + 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]); + + dpapimk->contents_len = contents_len; + + for (u32 i = 0; i < dpapimk->contents_len / 4; i++) + { + dpapimk->contents[i] = hex_to_u32 ((const u8 *) &contents_pos[i * 8]); + + dpapimk->contents[i] = byte_swap_32 (dpapimk->contents[i]); + } + + u32 SID_len = cipher_algo_pos - 1 - SID_pos; + + /* maximum size of SID supported */ + u8 *SID_unicode = (u8 *) hcmalloc (32 * 4); + memset (SID_unicode, 0, 32 * 4); + + for (u32 i = 0; i < SID_len; i += 1) + { + SID_unicode[i * 2] = SID_pos[i]; + } + + SID_unicode[(SID_len + 1) * 2] = 0x80; + + /* Specific to DPAPI: needs trailing '\0' while computing hash */ + dpapimk->SID_len = (SID_len + 1) * 2; + + memcpy ((u8 *) dpapimk->SID, SID_unicode, 32 * 4); + + for (u32 i = 0; i < 32; i++) + { + dpapimk->SID[i] = byte_swap_32 (dpapimk->SID[i]); + } + + digest[0] = dpapimk->iv[0]; + digest[1] = dpapimk->iv[1]; + digest[2] = dpapimk->iv[2]; + digest[3] = dpapimk->iv[3]; + + salt->salt_buf[0] = dpapimk->iv[0]; + salt->salt_buf[1] = dpapimk->iv[1]; + salt->salt_buf[2] = dpapimk->iv[2]; + salt->salt_buf[3] = dpapimk->iv[3]; + + salt->salt_len = 16; + + hcfree(SID_unicode); + + return (PARSER_OK); +} + int wpa_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig) { u32 *digest = (u32 *) hash_buf->digest; @@ -12447,7 +12607,7 @@ int krb5tgs_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN u8 *edata_ptr = (u8 *) krb5tgs->edata2; - krb5tgs->edata2_len = (data_len - 32) / 2 ; + krb5tgs->edata2_len = (data_len - 32) / 2; /* skip '$' */ for (u32 i = 16 * 2 + 1; i < (krb5tgs->edata2_len * 2) + (16 * 2 + 1); i += 2) @@ -15226,6 +15386,7 @@ char *strhashtype (const u32 hash_mode) case 15000: return ((char *) HT_15000); case 15100: return ((char *) HT_15100); case 15200: return ((char *) HT_15200); + case 15300: return ((char *) HT_15300); case 99999: return ((char *) HT_99999); } @@ -17952,11 +18113,11 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le u32 algorithm = (u32) keepass->algorithm; u32 keyfile_len = (u32) keepass->keyfile_len; - u32 *ptr_final_random_seed = (u32 *) keepass->final_random_seed ; - u32 *ptr_transf_random_seed = (u32 *) keepass->transf_random_seed ; - u32 *ptr_enc_iv = (u32 *) keepass->enc_iv ; - u32 *ptr_contents_hash = (u32 *) keepass->contents_hash ; - u32 *ptr_keyfile = (u32 *) keepass->keyfile ; + u32 *ptr_final_random_seed = (u32 *) keepass->final_random_seed; + u32 *ptr_transf_random_seed = (u32 *) keepass->transf_random_seed; + u32 *ptr_enc_iv = (u32 *) keepass->enc_iv; + u32 *ptr_contents_hash = (u32 *) keepass->contents_hash; + u32 *ptr_keyfile = (u32 *) keepass->keyfile; /* specific to version 2 */ u32 expected_bytes_len; @@ -18041,7 +18202,7 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le else if (version == 2) { expected_bytes_len = 8; - ptr_expected_bytes = (u32 *) keepass->expected_bytes ; + ptr_expected_bytes = (u32 *) keepass->expected_bytes; for (u32 i = 0; i < expected_bytes_len; i++, ptr_data += 8) sprintf (ptr_data, "%08x", ptr_expected_bytes[i]); @@ -18330,6 +18491,103 @@ 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); } + else if (hash_mode == 15300) + { + dpapimk_t *dpapimks = (dpapimk_t *) esalts_buf; + + dpapimk_t *dpapimk = &dpapimks[digest_cur]; + + u32 version = (u32) dpapimk->version; + u32 context = (u32) dpapimk->context; + u32 rounds = salt.salt_iter + 1; + u32 contents_len = (u32) dpapimk->contents_len; + u32 SID_len = (u32) dpapimk->SID_len; + u32 iv_len = 32; + + u8 cipher_algorithm[8] = { 0 }; + u8 hash_algorithm[8] = { 0 }; + u8 SID[512] = { 0 }; + u8* SID_tmp; + + u32 *ptr_SID = (u32 *) dpapimk->SID; + u32 *ptr_iv = (u32 *) dpapimk->iv; + u32 *ptr_contents = (u32 *) dpapimk->contents; + + u32 u32_iv[4]; + u8 iv[32 + 1]; + + /* convert back SID */ + + SID_tmp = (u8 *) hcmalloc ((SID_len + 1) * sizeof(u8)); + + for (u32 i = 0; i < (SID_len / 4) + 1; i++) + { + u8 hex[8] = { 0 }; + u32_to_hex_lower (byte_swap_32 (ptr_SID[i]), hex); + + for (u32 j = 0, k = 0; j < 8; j += 2, k++) + { + SID_tmp[i * 4 + k] = hex_to_u8 (&hex[j]); + } + } + /* overwrite trailing 0x80 */ + SID_tmp[SID_len] = 0; + + for (u32 i = 0, j = 0 ; j < SID_len ; i++, j += 2) + { + SID[i] = SID_tmp[j]; + } + + hcfree(SID_tmp); + + 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; + + u32 u32_contents[36]; + u8 contents[288 + 1]; + + 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); + } + + if (version == 1) + { + contents[208] = 0; + } + else + { + contents[288] = 0; + } + + if (contents_len == 288 && version == 2) + { + memcpy(cipher_algorithm, "aes256", strlen("aes256")); + memcpy(hash_algorithm, "sha512", strlen("sha512")); + } + else if (contents_len == 208 && version == 1) + { + memcpy(cipher_algorithm, "des3", strlen("des3")); + memcpy(hash_algorithm, "sha1", strlen("sha1")); + } + + snprintf (out_buf, out_len - 1, "%s%d*%d*%s*%s*%s*%d*%s*%d*%s", + SIGNATURE_DPAPIMK, + version, + context, + SID, + cipher_algorithm, + hash_algorithm, + rounds, + iv, + contents_len, + contents); + } else if (hash_mode == 99999) { char *ptr = (char *) digest_buf; @@ -22627,6 +22885,20 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->dgst_pos3 = 3; break; + 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; + hashconfig->kern_type = KERN_TYPE_DPAPIMK; + hashconfig->dgst_size = DGST_SIZE_4_4; + hashconfig->parse_func = dpapimk_parse_hash; + hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE; + hashconfig->dgst_pos0 = 0; + hashconfig->dgst_pos1 = 1; + hashconfig->dgst_pos2 = 2; + hashconfig->dgst_pos3 = 3; + break; + case 99999: hashconfig->hash_type = HASH_TYPE_PLAINTEXT; hashconfig->salt_type = SALT_TYPE_NONE; hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL; @@ -22763,6 +23035,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; + case 15300: hashconfig->esalt_size = sizeof (dpapimk_t); break; } // hook_salt_size @@ -22866,6 +23139,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; + case 15300: hashconfig->tmp_size = sizeof (dpapimk_tmp_t); break; }; // hook_size @@ -23270,6 +23544,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; + case 15300: ((dpapimk_t *) esalt)->version = 1; + break; } // special hook salt handling @@ -23464,6 +23740,8 @@ void hashconfig_benchmark_defaults (hashcat_ctx_t *hashcat_ctx, salt_t *salt, vo break; case 15200: salt->salt_iter = ROUNDS_MYWALLETV2; break; + case 15300: salt->salt_iter = ROUNDS_DPAPIMK; + break; } } diff --git a/src/usage.c b/src/usage.c index 7f44abfeb..f39f4fad8 100644 --- a/src/usage.c +++ b/src/usage.c @@ -241,6 +241,7 @@ static const char *USAGE_BIG[] = " 1000 | NTLM | Operating Systems", " 1100 | Domain Cached Credentials (DCC), MS Cache | Operating Systems", " 2100 | Domain Cached Credentials 2 (DCC2), MS Cache 2 | Operating Systems", + " 15300 | DPAPI masterkey file v1 and v2 | Operating Systems", " 12800 | MS-AzureSync PBKDF2-HMAC-SHA256 | Operating Systems", " 1500 | descrypt, DES (Unix), Traditional DES | Operating Systems", " 12400 | BSDiCrypt, Extended DES | Operating Systems", diff --git a/tools/test.pl b/tools/test.pl index 9a9565daf..98fa4e489 100755 --- a/tools/test.pl +++ b/tools/test.pl @@ -7,19 +7,19 @@ use strict; use warnings; -use Digest::MD4 qw (md4 md4_hex); -use Digest::MD5 qw (md5 md5_hex); -use Digest::SHA qw (sha1 sha256 sha384 sha512 sha1_hex sha224_hex sha256_hex sha384_hex sha512_hex); -use Digest::HMAC qw (hmac hmac_hex); -use Digest::Keccak qw (keccak_256_hex); -use Digest::BLAKE2 qw (blake2b_hex); -use Crypt::MySQL qw (password41); -use Digest::GOST qw (gost gost_hex); -use Digest::HMAC_MD5 qw (hmac_md5); -use Digest::CRC qw (crc32); +use Digest::MD4 qw (md4 md4_hex); +use Digest::MD5 qw (md5 md5_hex); +use Digest::SHA qw (sha1 sha256 sha384 sha512 sha1_hex sha224_hex sha256_hex sha384_hex sha512_hex hmac_sha1 hmac_sha512); +use Digest::HMAC qw (hmac hmac_hex); +use Digest::Keccak qw (keccak_256_hex); +use Digest::BLAKE2 qw (blake2b_hex); +use Crypt::MySQL qw (password41); +use Digest::GOST qw (gost gost_hex); +use Digest::HMAC_MD5 qw (hmac_md5); +use Digest::CRC qw (crc32); use Crypt::PBKDF2; use Crypt::DES; -use Crypt::ECB qw (encrypt); +use Crypt::ECB qw (encrypt); use Crypt::CBC; use Crypt::Eksblowfish::Bcrypt qw (bcrypt en_base64); use Crypt::Digest::RIPEMD160 qw (ripemd160_hex); @@ -47,11 +47,11 @@ my $hashcat = "./hashcat"; my $MAX_LEN = 55; -my @modes = (0, 10, 11, 12, 20, 21, 22, 23, 30, 40, 50, 60, 100, 101, 110, 111, 112, 120, 121, 122, 125, 130, 131, 132, 133, 140, 141, 150, 160, 200, 300, 400, 500, 600, 900, 1000, 1100, 1300, 1400, 1410, 1411, 1420, 1430, 1440, 1441, 1450, 1460, 1500, 1600, 1700, 1710, 1711, 1720, 1730, 1740, 1722, 1731, 1750, 1760, 1800, 2100, 2400, 2410, 2500, 2600, 2611, 2612, 2711, 2811, 3000, 3100, 3200, 3710, 3711, 3300, 3500, 3610, 3720, 3800, 3910, 4010, 4110, 4210, 4300, 4400, 4500, 4520, 4521, 4522, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7000, 7100, 7200, 7300, 7400, 7500, 7700, 7800, 7900, 8000, 8100, 8200, 8300, 8400, 8500, 8600, 8700, 8900, 9100, 9200, 9300, 9400, 9500, 9600, 9700, 9800, 9900, 10000, 10100, 10200, 10300, 10400, 10500, 10600, 10700, 10800, 10900, 11000, 11100, 11200, 11300, 11400, 11500, 11600, 11900, 12000, 12001, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100, 13200, 13300, 13400, 13500, 13600, 13800, 13900, 14000, 14100, 14400, 14700, 14800, 14900, 15000, 15100, 15200, 99999); +my @modes = (0, 10, 11, 12, 20, 21, 22, 23, 30, 40, 50, 60, 100, 101, 110, 111, 112, 120, 121, 122, 125, 130, 131, 132, 133, 140, 141, 150, 160, 200, 300, 400, 500, 600, 900, 1000, 1100, 1300, 1400, 1410, 1411, 1420, 1430, 1440, 1441, 1450, 1460, 1500, 1600, 1700, 1710, 1711, 1720, 1730, 1740, 1722, 1731, 1750, 1760, 1800, 2100, 2400, 2410, 2500, 2600, 2611, 2612, 2711, 2811, 3000, 3100, 3200, 3710, 3711, 3300, 3500, 3610, 3720, 3800, 3910, 4010, 4110, 4210, 4300, 4400, 4500, 4520, 4521, 4522, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7000, 7100, 7200, 7300, 7400, 7500, 7700, 7800, 7900, 8000, 8100, 8200, 8300, 8400, 8500, 8600, 8700, 8900, 9100, 9200, 9300, 9400, 9500, 9600, 9700, 9800, 9900, 10000, 10100, 10200, 10300, 10400, 10500, 10600, 10700, 10800, 10900, 11000, 11100, 11200, 11300, 11400, 11500, 11600, 11900, 12000, 12001, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100, 13200, 13300, 13400, 13500, 13600, 13800, 13900, 14000, 14100, 14400, 14700, 14800, 14900, 15000, 15100, 15200, 15300, 99999); -my %is_unicode = map { $_ => 1 } qw(30 40 130 131 132 133 140 141 1000 1100 1430 1440 1441 1730 1740 1731 5500 5600 8000 9400 9500 9600 9700 9800 11600 13500 13800); -my %less_fifteen = map { $_ => 1 } qw(500 1600 1800 2400 2410 3200 6300 7400 10500 10700); -my %allow_long_salt = map { $_ => 1 } qw(2500 4520 4521 5500 5600 7100 7200 7300 9400 9500 9600 9700 9800 10400 10500 10600 10700 1100 11000 11200 11300 11400 11600 12600 13500 13800 15000); +my %is_unicode = map { $_ => 1 } qw (30 40 130 131 132 133 140 141 1000 1100 1430 1440 1441 1730 1740 1731 5500 5600 8000 9400 9500 9600 9700 9800 11600 13500 13800); +my %less_fifteen = map { $_ => 1 } qw (500 1600 1800 2400 2410 3200 6300 7400 10500 10700); +my %allow_long_salt = map { $_ => 1 } qw (2500 4520 4521 5500 5600 7100 7200 7300 9400 9500 9600 9700 9800 10400 10500 10600 10700 1100 11000 11200 11300 11400 11600 12600 13500 13800 15000); my @lotus_magic_table = ( @@ -2621,6 +2621,60 @@ sub verify next unless (exists ($db->{$hash_in}) and (! defined ($db->{$hash_in}))); } + elsif ($mode == 15300) + { + ($hash_in, $word) = split ":", $line; + + next unless defined $hash_in; + next unless defined $word; + + my @tmp_data = split ('\$', $hash_in); + + my $signature = $tmp_data[1]; + + next unless ($signature eq 'DPAPImk'); + + my @data = split ('\*', $tmp_data[2]); + + next unless (scalar @data == 9); + + my $version = shift @data; + + next unless ($version == 1 || $version == 2); + + my $context = shift @data; + + my $SID = shift @data; + + my $cipher_algorithm = shift @data; + + my $hash_algorithm = shift @data; + + my $iteration = shift @data; + + my $iv = shift @data; + + my $cipher_len = shift @data; + + my $cipher = shift @data; + + next unless (length ($cipher) == $cipher_len); + + if ($version == 1) + { + next unless ($cipher_len == 208); + } + elsif ($version == 2) + { + next unless ($cipher_len == 288); + } + + $salt = substr ($hash_in, length ('$DPAPImk$')); + + $param = $cipher; + + next unless (exists ($db->{$hash_in}) and (! defined ($db->{$hash_in}))); + } else { print "ERROR: hash mode is not supported\n"; @@ -2996,6 +3050,14 @@ sub verify $hash_out = $hash_in; } + elsif ($mode == 15300) + { + $hash_out = gen_hash ($mode, $word, $salt, $iter, $param); + + $len = length $hash_out; + + return unless (substr ($line, 0, $len) eq $hash_out); + } else { $hash_out = gen_hash ($mode, $word, $salt, $iter); @@ -3510,6 +3572,12 @@ sub passthrough { $tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 32)); } + elsif ($mode == 15300) + { + $salt_buf = get_random_dpapimk_salt (); + + $tmp_hash = gen_hash ($mode, $word_buf, $salt_buf); + } else { print "ERROR: Unsupported hash type\n"; @@ -4450,6 +4518,20 @@ sub single } } } + elsif ($mode == 15300) + { + for (my $i = 1; $i < 16; $i++) + { + if ($len != 0) + { + rnd ($mode, $len, 16); + } + else + { + rnd ($mode, $i, 16); + } + } + } } } @@ -4780,7 +4862,7 @@ sub gen_hash } elsif ($mode == 600) { - $hash_buf = lc blake2b_hex($word_buf); + $hash_buf = lc blake2b_hex ($word_buf); $tmp_hash = sprintf ("\$BLAKE2\$" . $hash_buf); } elsif ($mode == 900) @@ -8153,6 +8235,266 @@ END_CODE $tmp_hash = sprintf ("\$blockchain\$v2\$%d\$%s\$%s", $iterations, length ($salt_buf . $encrypted) / 2, $salt_buf . $encrypted); } + elsif ($mode == 15300) + { + my @salt_arr = split ('\*', $salt_buf); + + my $version = $salt_arr[0]; + + my $context = $salt_arr[1]; + + my $SID = $salt_arr[2]; + + my $cipher_algorithm = $salt_arr[3]; + + my $hash_algorithm = $salt_arr[4]; + + my $iterations = $salt_arr[5]; + + my $salt = pack ("H*", $salt_arr[6]); + + my $cipher_len = $salt_arr[7]; + + my $cipher; + + # intermediate values + + my $user_hash; + my $user_derivationKey; + my $encKey; + my $expected_hmac; + my $cleartext; + + if ($context == 1) + { + $user_hash = sha1 (encode ("UTF-16LE", $word_buf)); + } + elsif ($context == 2) + { + $user_hash = md4 (encode ("UTF-16LE", $word_buf)); + } + + $user_derivationKey = hmac_sha1 (encode ("UTF-16LE", $SID . "\x00"), $user_hash); + + my $hmacSalt = randbytes (16); + my $last_key = randbytes (64); + + if ($version == 1) + { + $encKey = hmac_sha1 ($hmacSalt, $user_derivationKey); + $expected_hmac = hmac_sha1 ($last_key, $encKey); + + # need padding because keyLen is 24 and hashLen 20 + $expected_hmac = $expected_hmac . randbytes (4); + } + elsif ($version == 2) + { + $encKey = hmac_sha512 ($hmacSalt, $user_derivationKey); + $expected_hmac = hmac_sha512 ($last_key, $encKey); + } + + $cleartext = $hmacSalt . $expected_hmac . $last_key; + + my $derived_key; + my $key; + my $iv; + + my $pbkdf2; + + if ($version == 1) + { + $derived_key = dpapi_pbkdf2 ($user_derivationKey, $salt, $iterations, 32, \&hmac_sha1); + } + elsif ($version == 2) + { + $derived_key = dpapi_pbkdf2 ($user_derivationKey, $salt, $iterations, 48, \&hmac_sha512); + } + + if (defined $additional_param) + { + $cipher = pack ("H*", $additional_param); + my $computed_hmac = ""; + + if ($version == 1) + { + $key = substr ($derived_key, 0, 24); + $iv = substr ($derived_key, 24, 8); + + my $p1 = Crypt::ECB->new ({ + key => substr ($key, 0, 8), + cipher => "DES", + literal_key => 1, + header => "none", + keysize => 8, + padding => "null", + }); + + my $p2 = Crypt::ECB->new ({ + key => substr ($key, 8, 8), + cipher => "DES", + literal_key => 1, + header => "none", + keysize => 8, + padding => "null", + }); + + my $p3 = Crypt::ECB->new ({ + key => substr ($key, 16, 8), + cipher => "DES", + literal_key => 1, + header => "none", + keysize => 8, + padding => "null", + }); + + # let's compute a 3DES-EDE-CBC decryption + + my $out1; + my $out2; + my $out3; + my $expected_cleartext = ""; + + # size of cipherlen is 104 bytes + for (my $k = 0; $k < 13; $k++) + { + $out1 = $p3->decrypt (substr ($cipher, $k * 8, 8)); + $out2 = $p2->encrypt ($out1); + $out3 = $p1->decrypt ($out2); + + $expected_cleartext .= substr ($out3, 0, 8) ^ $iv; + + $iv = substr ($cipher, $k * 8, 8); + } + + $last_key = substr ($expected_cleartext, length ($expected_cleartext) - 64, 64); + $hmacSalt = substr ($expected_cleartext, 0, 16); + $expected_hmac = substr ($expected_cleartext, 16, 20); + + $encKey = hmac_sha1 ($hmacSalt, $user_derivationKey); + $computed_hmac = hmac_sha1 ($last_key, $encKey); + + $cleartext = $expected_cleartext; + + if (unpack ("H*", $expected_hmac) ne unpack ("H*", $computed_hmac)) + { + $cleartext = "0" x 104; + } + } + elsif ($version == 2) + { + $key = substr ($derived_key, 0, 32); + $iv = substr ($derived_key, 32, 16); + + my $aes = Crypt::CBC->new ({ + key => $key, + cipher => "Crypt::Rijndael", + iv => $iv, + literal_key => 1, + header => "none", + keysize => 32, + padding => "null", + }); + + my $expected_cleartext = $aes->decrypt ($cipher); + + $last_key = substr ($expected_cleartext, length ($expected_cleartext) - 64, 64); + $hmacSalt = substr ($expected_cleartext, 0, 16); + $expected_hmac = substr ($expected_cleartext, 16, 64); + + $encKey = hmac_sha512 ($hmacSalt, $user_derivationKey); + $computed_hmac = hmac_sha512 ($last_key, $encKey); + + $cleartext = $expected_cleartext; + + if (unpack ("H*", $expected_hmac) ne unpack ("H*", $computed_hmac)) + { + $cleartext = "0" x 144; + } + } + } + + if ($version == 1) + { + $key = substr ($derived_key, 0, 24); + $iv = substr ($derived_key, 24, 8); + + my $p1 = Crypt::ECB->new ({ + key => substr ($key, 0, 8), + cipher => "DES", + literal_key => 1, + header => "none", + keysize => 8, + padding => "null", + }); + + my $p2 = Crypt::ECB->new ({ + key => substr ($key, 8, 8), + cipher => "DES", + literal_key => 1, + header => "none", + keysize => 8, + padding => "null", + }); + + my $p3 = Crypt::ECB->new ({ + key => substr ($key, 16, 8), + cipher => "DES", + literal_key => 1, + header => "none", + keysize => 8, + padding => "null", + }); + + # let's compute a 3DES-EDE-CBC encryption + + # compute first block + my $out1 = $p1->encrypt (substr ($cleartext, 0, 8) ^ $iv); + my $out2 = $p2->decrypt ($out1); + my $out3 = $p3->encrypt ($out2); + + $cipher = substr ($out3, 0, 8); + + # size of cipherlen is 104 bytes + for (my $k = 1; $k < 13; $k++) + { + $iv = $out3; + + $out1 = $p1->encrypt (substr ($cleartext, $k * 8, 8) ^ $iv); + $out2 = $p2->decrypt ($out1); + $out3 = $p3->encrypt ($out2); + + $cipher .= substr ($out3, 0, 8); + } + } + else + { + $key = substr ($derived_key, 0, 32); + $iv = substr ($derived_key, 32, 16); + + my $aes = Crypt::CBC->new ({ + key => $key, + cipher => "Crypt::Rijndael", + iv => $iv, + literal_key => 1, + header => "none", + keysize => 32, + padding => "null", + }); + + $cipher = $aes->encrypt ($cleartext); + } + + $tmp_hash = sprintf ('$DPAPImk$%d*%d*%s*%s*%s*%d*%s*%d*%s', + $version, + $context, + $SID, + $cipher_algorithm, + $hash_algorithm, + $iterations, + unpack ("H*", $salt), + $cipher_len, + unpack ("H*", $cipher)); + } elsif ($mode == 99999) { $tmp_hash = sprintf ("%s", $word_buf); @@ -8161,6 +8503,30 @@ END_CODE return ($tmp_hash); } +#Thanks to Jochen Hoenicke +# (one of the authors of Palm Keyring) +# for these next two subs. +sub dpapi_pbkdf2 +{ + my ($password, $salt, $iter, $keylen, $prf) = @_; + my ($k, $t, $u, $ui, $i); + $t = ""; + for ($k = 1; length ($t) < $keylen; $k++) + { + $u = $ui = &$prf ($salt.pack ('N', $k), $password); + for ($i = 1; $i < $iter; $i++) + { + # modification to fit Microsoft + # weird pbkdf2 implementation... + $ui = &$prf ($u, $password); + $u ^= $ui; + } + $t .= $u; + } + return substr ($t, 0, $keylen); +} + + sub rnd { my $mode = shift; @@ -8267,6 +8633,10 @@ sub rnd { $salt_buf = get_pstoken_salt (); } + elsif ($mode == 15300) + { + $salt_buf = get_random_dpapimk_salt (); + } else { my @salt_arr; @@ -8562,9 +8932,9 @@ sub pseudo_base64 my $md5 = shift; my $s64 = ""; for my $i (0..3) { - my $v = unpack "V", substr($md5, $i*4, 4); + my $v = unpack "V", substr ($md5, $i*4, 4); for (1..4) { - $s64 .= substr($itoa64, $v & 0x3f, 1); + $s64 .= substr ($itoa64, $v & 0x3f, 1); $v >>= 6; } } @@ -8606,31 +8976,31 @@ sub oracle_hash { my ($username, $password) = @_; - my $userpass = pack('n*', unpack('C*', uc($username.$password))); - $userpass .= pack('C', 0) while (length($userpass) % 8); + my $userpass = pack ('n*', unpack ('C*', uc ($username.$password))); + $userpass .= pack ('C', 0) while (length ($userpass) % 8); - my $key = pack('H*', "0123456789ABCDEF"); - my $iv = pack('H*', "0000000000000000"); + my $key = pack ('H*', "0123456789ABCDEF"); + my $iv = pack ('H*', "0000000000000000"); - my $c = new Crypt::CBC( + my $c = new Crypt::CBC ( -literal_key => 1, -cipher => "DES", -key => $key, -iv => $iv, -header => "none" ); - my $key2 = substr($c->encrypt($userpass), length($userpass)-8, 8); + my $key2 = substr ($c->encrypt ($userpass), length ($userpass)-8, 8); - my $c2 = new Crypt::CBC( + my $c2 = new Crypt::CBC ( -literal_key => 1, -cipher => "DES", -key => $key2, -iv => $iv, -header => "none" ); - my $hash = substr($c2->encrypt($userpass), length($userpass)-8, 8); + my $hash = substr ($c2->encrypt ($userpass), length ($userpass)-8, 8); - return uc(unpack('H*', $hash)); + return uc (unpack ('H*', $hash)); } sub androidpin_hash @@ -9291,13 +9661,13 @@ sub setup_des_key $key = $key_56[0]; - $key .= chr(((ord($key_56[0]) << 7) | (ord($key_56[1]) >> 1)) & 255); - $key .= chr(((ord($key_56[1]) << 6) | (ord($key_56[2]) >> 2)) & 255); - $key .= chr(((ord($key_56[2]) << 5) | (ord($key_56[3]) >> 3)) & 255); - $key .= chr(((ord($key_56[3]) << 4) | (ord($key_56[4]) >> 4)) & 255); - $key .= chr(((ord($key_56[4]) << 3) | (ord($key_56[5]) >> 5)) & 255); - $key .= chr(((ord($key_56[5]) << 2) | (ord($key_56[6]) >> 6)) & 255); - $key .= chr(( ord($key_56[6]) << 1) & 255); + $key .= chr (((ord ($key_56[0]) << 7) | (ord ($key_56[1]) >> 1)) & 255); + $key .= chr (((ord ($key_56[1]) << 6) | (ord ($key_56[2]) >> 2)) & 255); + $key .= chr (((ord ($key_56[2]) << 5) | (ord ($key_56[3]) >> 3)) & 255); + $key .= chr (((ord ($key_56[3]) << 4) | (ord ($key_56[4]) >> 4)) & 255); + $key .= chr (((ord ($key_56[4]) << 3) | (ord ($key_56[5]) >> 5)) & 255); + $key .= chr (((ord ($key_56[5]) << 2) | (ord ($key_56[6]) >> 6)) & 255); + $key .= chr (( ord ($key_56[6]) << 1) & 255); return $key; } @@ -9707,6 +10077,64 @@ sub get_random_dnssec_salt return $salt_buf; } +sub get_random_dpapimk_salt +{ + my $salt_buf = ""; + + my $version = get_random_num (1, 3); + + my $context = get_random_num (1, 3); + + my $cipher_algo = ""; + + my $hash_algo = ""; + + my $iterations; + + my $SID = sprintf ('S-15-21-%d-%d-%d-%d', + get_random_num (400000000,490000000), + get_random_num (400000000,490000000), + get_random_num (400000000,490000000), + get_random_num (1000,1999)); + + my $cipher_len = 0; + + if ($version == 1) + { + $iterations = get_random_num (4000, 24000); + + $cipher_algo = "des3"; + + $hash_algo = "sha1"; + + $cipher_len = 208; + } + elsif ($version == 2) + { + $iterations = get_random_num (8000, 17000); + + $cipher_algo = "aes256"; + + $hash_algo = "sha512"; + + $cipher_len = 288; + } + + my $iv = randbytes (16); + $iv = unpack ("H*", $iv); + + $salt_buf = $version . '*' . + $context . '*' . + $SID . '*' . + $cipher_algo . '*' . + $hash_algo . '*' . + $iterations . '*' . + $iv . '*' . + $cipher_len . '*'; + + return $salt_buf; +} + sub md5bit { my $digest = shift; diff --git a/tools/test.sh b/tools/test.sh index ea7893239..236e5e7ec 100755 --- a/tools/test.sh +++ b/tools/test.sh @@ -9,7 +9,7 @@ TDIR="$( cd "$( dirname "${BASH_SOURCE[0]}" )" && pwd )" # missing hash types: 5200,6251,6261,6271,6281 -HASH_TYPES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 125 130 131 132 133 140 141 150 160 200 300 400 500 600 900 1000 1100 1300 1400 1410 1411 1420 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 3910 4010 4110 4300 4400 4500 4520 4521 4522 4700 4800 4900 5000 5100 5300 5400 5500 5600 5700 5800 6000 6100 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6300 6400 6500 6600 6700 6800 6900 7000 7100 7200 7300 7400 7500 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8900 9100 9200 9300 9400 9500 9600 9700 9800 9900 10000 10100 10200 10300 10400 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11900 12000 12001 12100 12200 12300 12400 12600 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 14000 14100 14400 14600 14700 14800 14900 15000 15100 15200 99999" +HASH_TYPES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 125 130 131 132 133 140 141 150 160 200 300 400 500 600 900 1000 1100 1300 1400 1410 1411 1420 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 3910 4010 4110 4300 4400 4500 4520 4521 4522 4700 4800 4900 5000 5100 5300 5400 5500 5600 5700 5800 6000 6100 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6300 6400 6500 6600 6700 6800 6900 7000 7100 7200 7300 7400 7500 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8900 9100 9200 9300 9400 9500 9600 9700 9800 9900 10000 10100 10200 10300 10400 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11900 12000 12001 12100 12200 12300 12400 12600 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 14000 14100 14400 14600 14700 14800 14900 15000 15100 15200 15300 99999" #ATTACK_MODES="0 1 3 6 7" ATTACK_MODES="0 1 3 7" @@ -22,7 +22,7 @@ HASHFILE_ONLY="2500" NEVER_CRACK="11600 14900" -SLOW_ALGOS="400 500 501 1600 1800 2100 2500 3200 5200 5800 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6251 6261 6271 6281 6300 6400 6500 6600 6700 6800 7100 7200 7400 7900 8200 8800 8900 9000 9100 9200 9300 9400 9500 9600 10000 10300 10500 10700 10900 11300 11600 11900 12000 12001 12100 12200 12300 12400 12500 12700 12800 12900 13000 13200 13400 13600 14600 14700 14800 15100 15200" +SLOW_ALGOS="400 500 501 1600 1800 2100 2500 3200 5200 5800 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6251 6261 6271 6281 6300 6400 6500 6600 6700 6800 7100 7200 7400 7900 8200 8800 8900 9000 9100 9200 9300 9400 9500 9600 10000 10300 10500 10700 10900 11300 11600 11900 12000 12001 12100 12200 12300 12400 12500 12700 12800 12900 13000 13200 13400 13600 14600 14700 14800 15100 15200 15300" OPTS="--quiet --force --potfile-disable --runtime 400 --gpu-temp-disable --weak-hash-threshold=0"