mirror of
https://github.com/hashcat/hashcat
synced 2024-11-02 20:39:22 +01:00
Two new modes added:
* Added new hash-mode 14000 = DES (PT = $salt, key = $pass) * Added new hash-mode 14100 = 3DES (PT = $salt, key = $pass) This patch initial was created out of this PR: https://github.com/hashcat/hashcat/pull/452 which is now no longer required I've replace the -a 3 kernel with a bitsliced version which is around 10 times faster I've also added test.sh and test.pl patches, but some of the test scripts need to be fixed
This commit is contained in:
jsteube
parent
18d6a7a5d5
commit
71a8459d85
@ -128,6 +128,12 @@
|
||||
#if KERN_TYPE == 13723
|
||||
#undef _unroll
|
||||
#endif
|
||||
#if KERN_TYPE == 14000
|
||||
#undef _unroll
|
||||
#endif
|
||||
#if KERN_TYPE == 14100
|
||||
#undef _unroll
|
||||
#endif
|
||||
|
||||
#endif
|
||||
#endif
|
||||
|
||||
741
OpenCL/m14000_a0.cl
Normal file
741
OpenCL/m14000_a0.cl
Normal file
@ -0,0 +1,741 @@
|
||||
/**
|
||||
* Authors.....: Jens Steube <jens.steube@gmail.com>
|
||||
* Gabriele Gristina <matrix@hashcat.net>
|
||||
* magnum <john.magnum@hushmail.com>
|
||||
* Frans Lategan <frans.lategan+hashcat@gmail.com>
|
||||
*
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
#define _DES_
|
||||
|
||||
#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_rp.h"
|
||||
#include "inc_rp.cl"
|
||||
#include "inc_simd.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 u32 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 u32 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
|
||||
}
|
||||
};
|
||||
|
||||
#if VECT_SIZE == 1
|
||||
#define BOX(i,n,S) (S)[(n)][(i)]
|
||||
#elif VECT_SIZE == 2
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
|
||||
#elif VECT_SIZE == 4
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
|
||||
#elif VECT_SIZE == 8
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7])
|
||||
#elif VECT_SIZE == 16
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
|
||||
#endif
|
||||
|
||||
#if VECT_SIZE == 1
|
||||
#define BOX1(i,S) (S)[(i)]
|
||||
#elif VECT_SIZE == 2
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1])
|
||||
#elif VECT_SIZE == 4
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3])
|
||||
#elif VECT_SIZE == 8
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7])
|
||||
#elif VECT_SIZE == 16
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7], (S)[(i).s8], (S)[(i).s9], (S)[(i).sa], (S)[(i).sb], (S)[(i).sc], (S)[(i).sd], (S)[(i).se], (S)[(i).sf])
|
||||
#endif
|
||||
|
||||
void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
|
||||
{
|
||||
u32x r = rotl32 (data[0], 3u);
|
||||
u32x l = rotl32 (data[1], 3u);
|
||||
|
||||
u32x tt;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i += 2)
|
||||
{
|
||||
u32x u;
|
||||
u32x t;
|
||||
|
||||
u = Kc[i + 0] ^ r;
|
||||
t = Kd[i + 0] ^ rotl32 (r, 28u);
|
||||
|
||||
l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
|
||||
| BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
|
||||
| BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
|
||||
| BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
|
||||
| BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
|
||||
| BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
|
||||
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
|
||||
|
||||
u = Kc[i + 1] ^ l;
|
||||
t = Kd[i + 1] ^ rotl32 (l, 28u);
|
||||
|
||||
r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
|
||||
| BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
|
||||
| BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
|
||||
| BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
|
||||
| BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
|
||||
| BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
|
||||
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
|
||||
}
|
||||
|
||||
iv[0] = rotl32 (l, 29u);
|
||||
iv[1] = rotl32 (r, 29u);
|
||||
}
|
||||
|
||||
void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
|
||||
{
|
||||
u32x tt;
|
||||
|
||||
PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
|
||||
HPERM_OP (c, tt, 2, 0xcccc0000);
|
||||
HPERM_OP (d, tt, 2, 0xcccc0000);
|
||||
PERM_OP (d, c, tt, 1, 0x55555555);
|
||||
PERM_OP (c, d, tt, 8, 0x00ff00ff);
|
||||
PERM_OP (d, c, tt, 1, 0x55555555);
|
||||
|
||||
d = ((d & 0x000000ff) << 16)
|
||||
| ((d & 0x0000ff00) << 0)
|
||||
| ((d & 0x00ff0000) >> 16)
|
||||
| ((c & 0xf0000000) >> 4);
|
||||
|
||||
c = c & 0x0fffffff;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i++)
|
||||
{
|
||||
if ((i < 2) || (i == 8) || (i == 15))
|
||||
{
|
||||
c = ((c >> 1) | (c << 27));
|
||||
d = ((d >> 1) | (d << 27));
|
||||
}
|
||||
else
|
||||
{
|
||||
c = ((c >> 2) | (c << 26));
|
||||
d = ((d >> 2) | (d << 26));
|
||||
}
|
||||
|
||||
c = c & 0x0fffffff;
|
||||
d = d & 0x0fffffff;
|
||||
|
||||
const u32x c00 = (c >> 0) & 0x0000003f;
|
||||
const u32x c06 = (c >> 6) & 0x00383003;
|
||||
const u32x c07 = (c >> 7) & 0x0000003c;
|
||||
const u32x c13 = (c >> 13) & 0x0000060f;
|
||||
const u32x c20 = (c >> 20) & 0x00000001;
|
||||
|
||||
u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
|
||||
| BOX (((c06 >> 0) & 0xff)
|
||||
|((c07 >> 0) & 0xff), 1, s_skb)
|
||||
| BOX (((c13 >> 0) & 0xff)
|
||||
|((c06 >> 8) & 0xff), 2, s_skb)
|
||||
| BOX (((c20 >> 0) & 0xff)
|
||||
|((c13 >> 8) & 0xff)
|
||||
|((c06 >> 16) & 0xff), 3, s_skb);
|
||||
|
||||
const u32x d00 = (d >> 0) & 0x00003c3f;
|
||||
const u32x d07 = (d >> 7) & 0x00003f03;
|
||||
const u32x d21 = (d >> 21) & 0x0000000f;
|
||||
const u32x d22 = (d >> 22) & 0x00000030;
|
||||
|
||||
u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
|
||||
| BOX (((d07 >> 0) & 0xff)
|
||||
|((d00 >> 8) & 0xff), 5, s_skb)
|
||||
| BOX (((d07 >> 8) & 0xff), 6, s_skb)
|
||||
| BOX (((d21 >> 0) & 0xff)
|
||||
|((d22 >> 0) & 0xff), 7, s_skb);
|
||||
|
||||
Kc[i] = ((t << 16) | (s & 0x0000ffff));
|
||||
Kd[i] = ((s >> 16) | (t & 0xffff0000));
|
||||
|
||||
Kc[i] = rotl32 (Kc[i], 2u);
|
||||
Kd[i] = rotl32 (Kd[i], 2u);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m14000_m04 (__global pw_t *pws, __global kernel_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_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);
|
||||
const u32 lid = get_local_id (0);
|
||||
const u32 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* 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;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
u32 pw_buf0[4];
|
||||
u32 pw_buf1[4];
|
||||
|
||||
pw_buf0[0] = pws[gid].i[ 0];
|
||||
pw_buf0[1] = pws[gid].i[ 1];
|
||||
pw_buf0[2] = 0;
|
||||
pw_buf0[3] = 0;
|
||||
pw_buf1[0] = 0;
|
||||
pw_buf1[1] = 0;
|
||||
pw_buf1[2] = 0;
|
||||
pw_buf1[3] = 0;
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
/**
|
||||
* salt
|
||||
*/
|
||||
|
||||
u32 salt_buf0[2];
|
||||
|
||||
salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
|
||||
salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
|
||||
|
||||
/**
|
||||
* main
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
|
||||
{
|
||||
u32x w0[4] = { 0 };
|
||||
u32x w1[4] = { 0 };
|
||||
u32x w2[4] = { 0 };
|
||||
u32x w3[4] = { 0 };
|
||||
|
||||
apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
|
||||
|
||||
const u32x c = w0[0];
|
||||
const u32x d = w0[1];
|
||||
|
||||
u32x Kc[16];
|
||||
u32x Kd[16];
|
||||
|
||||
_des_crypt_keysetup (c, d, Kc, Kd, s_skb);
|
||||
|
||||
u32x data[2];
|
||||
|
||||
data[0] = salt_buf0[0];
|
||||
data[1] = salt_buf0[1];
|
||||
|
||||
u32x iv[2];
|
||||
|
||||
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
|
||||
|
||||
u32x z = 0;
|
||||
|
||||
COMPARE_M_SIMD (iv[0], iv[1], z, z);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m14000_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m14000_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m14000_s04 (__global pw_t *pws, __global kernel_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_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);
|
||||
const u32 lid = get_local_id (0);
|
||||
const u32 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* 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;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
u32 pw_buf0[4];
|
||||
u32 pw_buf1[4];
|
||||
|
||||
pw_buf0[0] = pws[gid].i[ 0];
|
||||
pw_buf0[1] = pws[gid].i[ 1];
|
||||
pw_buf0[2] = 0;
|
||||
pw_buf0[3] = 0;
|
||||
pw_buf1[0] = 0;
|
||||
pw_buf1[1] = 0;
|
||||
pw_buf1[2] = 0;
|
||||
pw_buf1[3] = 0;
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
/**
|
||||
* salt
|
||||
*/
|
||||
|
||||
u32 salt_buf0[2];
|
||||
|
||||
salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
|
||||
salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
|
||||
|
||||
/**
|
||||
* digest
|
||||
*/
|
||||
|
||||
const u32 search[4] =
|
||||
{
|
||||
digests_buf[digests_offset].digest_buf[DGST_R0],
|
||||
digests_buf[digests_offset].digest_buf[DGST_R1],
|
||||
0,
|
||||
0
|
||||
};
|
||||
|
||||
/**
|
||||
* main
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
|
||||
{
|
||||
u32x w0[4] = { 0 };
|
||||
u32x w1[4] = { 0 };
|
||||
u32x w2[4] = { 0 };
|
||||
u32x w3[4] = { 0 };
|
||||
|
||||
apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
|
||||
|
||||
const u32x c = w0[0];
|
||||
const u32x d = w0[1];
|
||||
|
||||
u32x Kc[16];
|
||||
u32x Kd[16];
|
||||
|
||||
_des_crypt_keysetup (c, d, Kc, Kd, s_skb);
|
||||
|
||||
u32x data[2];
|
||||
|
||||
data[0] = salt_buf0[0];
|
||||
data[1] = salt_buf0[1];
|
||||
|
||||
u32x iv[2];
|
||||
|
||||
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
|
||||
|
||||
u32x z = 0;
|
||||
|
||||
COMPARE_S_SIMD (iv[0], iv[1], z, z);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m14000_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m14000_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
}
|
||||
817
OpenCL/m14000_a1.cl
Normal file
817
OpenCL/m14000_a1.cl
Normal file
@ -0,0 +1,817 @@
|
||||
/**
|
||||
* Authors.....: Jens Steube <jens.steube@gmail.com>
|
||||
* Gabriele Gristina <matrix@hashcat.net>
|
||||
* Frans Lategan <frans.lategan+hashcat@gmail.com>
|
||||
*
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
#define _DES_
|
||||
|
||||
#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"
|
||||
|
||||
#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 u32 c_SPtrans[8][64] =
|
||||
{
|
||||
{
|
||||
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,
|
||||
},
|
||||
{
|
||||
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,
|
||||
},
|
||||
{
|
||||
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,
|
||||
},
|
||||
{
|
||||
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,
|
||||
},
|
||||
{
|
||||
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,
|
||||
},
|
||||
{
|
||||
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,
|
||||
},
|
||||
{
|
||||
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,
|
||||
},
|
||||
{
|
||||
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 u32 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
|
||||
}
|
||||
};
|
||||
|
||||
#if VECT_SIZE == 1
|
||||
#define BOX(i,n,S) (S)[(n)][(i)]
|
||||
#elif VECT_SIZE == 2
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
|
||||
#elif VECT_SIZE == 4
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
|
||||
#elif VECT_SIZE == 8
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7])
|
||||
#elif VECT_SIZE == 16
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
|
||||
#endif
|
||||
|
||||
#if VECT_SIZE == 1
|
||||
#define BOX1(i,S) (S)[(i)]
|
||||
#elif VECT_SIZE == 2
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1])
|
||||
#elif VECT_SIZE == 4
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3])
|
||||
#elif VECT_SIZE == 8
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7])
|
||||
#elif VECT_SIZE == 16
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7], (S)[(i).s8], (S)[(i).s9], (S)[(i).sa], (S)[(i).sb], (S)[(i).sc], (S)[(i).sd], (S)[(i).se], (S)[(i).sf])
|
||||
#endif
|
||||
|
||||
void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
|
||||
{
|
||||
u32x r = rotl32 (data[0], 3u);
|
||||
u32x l = rotl32 (data[1], 3u);
|
||||
|
||||
u32x tt;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i += 2)
|
||||
{
|
||||
u32x u;
|
||||
u32x t;
|
||||
|
||||
u = Kc[i + 0] ^ r;
|
||||
t = Kd[i + 0] ^ rotl32 (r, 28u);
|
||||
|
||||
l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
|
||||
| BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
|
||||
| BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
|
||||
| BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
|
||||
| BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
|
||||
| BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
|
||||
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
|
||||
|
||||
u = Kc[i + 1] ^ l;
|
||||
t = Kd[i + 1] ^ rotl32 (l, 28u);
|
||||
|
||||
r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
|
||||
| BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
|
||||
| BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
|
||||
| BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
|
||||
| BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
|
||||
| BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
|
||||
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
|
||||
}
|
||||
|
||||
iv[0] = rotl32 (l, 29u);
|
||||
iv[1] = rotl32 (r, 29u);
|
||||
}
|
||||
|
||||
void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
|
||||
{
|
||||
u32x tt;
|
||||
|
||||
PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
|
||||
HPERM_OP (c, tt, 2, 0xcccc0000);
|
||||
HPERM_OP (d, tt, 2, 0xcccc0000);
|
||||
PERM_OP (d, c, tt, 1, 0x55555555);
|
||||
PERM_OP (c, d, tt, 8, 0x00ff00ff);
|
||||
PERM_OP (d, c, tt, 1, 0x55555555);
|
||||
|
||||
d = ((d & 0x000000ff) << 16)
|
||||
| ((d & 0x0000ff00) << 0)
|
||||
| ((d & 0x00ff0000) >> 16)
|
||||
| ((c & 0xf0000000) >> 4);
|
||||
|
||||
c = c & 0x0fffffff;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i++)
|
||||
{
|
||||
if ((i < 2) || (i == 8) || (i == 15))
|
||||
{
|
||||
c = ((c >> 1) | (c << 27));
|
||||
d = ((d >> 1) | (d << 27));
|
||||
}
|
||||
else
|
||||
{
|
||||
c = ((c >> 2) | (c << 26));
|
||||
d = ((d >> 2) | (d << 26));
|
||||
}
|
||||
|
||||
c = c & 0x0fffffff;
|
||||
d = d & 0x0fffffff;
|
||||
|
||||
const u32x c00 = (c >> 0) & 0x0000003f;
|
||||
const u32x c06 = (c >> 6) & 0x00383003;
|
||||
const u32x c07 = (c >> 7) & 0x0000003c;
|
||||
const u32x c13 = (c >> 13) & 0x0000060f;
|
||||
const u32x c20 = (c >> 20) & 0x00000001;
|
||||
|
||||
u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
|
||||
| BOX (((c06 >> 0) & 0xff)
|
||||
|((c07 >> 0) & 0xff), 1, s_skb)
|
||||
| BOX (((c13 >> 0) & 0xff)
|
||||
|((c06 >> 8) & 0xff), 2, s_skb)
|
||||
| BOX (((c20 >> 0) & 0xff)
|
||||
|((c13 >> 8) & 0xff)
|
||||
|((c06 >> 16) & 0xff), 3, s_skb);
|
||||
|
||||
const u32x d00 = (d >> 0) & 0x00003c3f;
|
||||
const u32x d07 = (d >> 7) & 0x00003f03;
|
||||
const u32x d21 = (d >> 21) & 0x0000000f;
|
||||
const u32x d22 = (d >> 22) & 0x00000030;
|
||||
|
||||
u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
|
||||
| BOX (((d07 >> 0) & 0xff)
|
||||
|((d00 >> 8) & 0xff), 5, s_skb)
|
||||
| BOX (((d07 >> 8) & 0xff), 6, s_skb)
|
||||
| BOX (((d21 >> 0) & 0xff)
|
||||
|((d22 >> 0) & 0xff), 7, s_skb);
|
||||
|
||||
Kc[i] = ((t << 16) | (s & 0x0000ffff));
|
||||
Kd[i] = ((s >> 16) | (t & 0xffff0000));
|
||||
|
||||
Kc[i] = rotl32 (Kc[i], 2u);
|
||||
Kd[i] = rotl32 (Kd[i], 2u);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m14000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_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);
|
||||
const u32 lid = get_local_id (0);
|
||||
const u32 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* 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;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
u32 pw_buf0[4];
|
||||
u32 pw_buf1[4];
|
||||
|
||||
pw_buf0[0] = pws[gid].i[ 0];
|
||||
pw_buf0[1] = pws[gid].i[ 1];
|
||||
pw_buf0[2] = 0;
|
||||
pw_buf0[3] = 0;
|
||||
pw_buf1[0] = 0;
|
||||
pw_buf1[1] = 0;
|
||||
pw_buf1[2] = 0;
|
||||
pw_buf1[3] = 0;
|
||||
|
||||
const u32 pw_l_len = pws[gid].pw_len;
|
||||
|
||||
/**
|
||||
* salt
|
||||
*/
|
||||
|
||||
u32 salt_buf0[2];
|
||||
|
||||
salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
|
||||
salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
|
||||
{
|
||||
const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
|
||||
|
||||
const u32x pw_len = pw_l_len + pw_r_len;
|
||||
|
||||
/**
|
||||
* concat password candidate
|
||||
*/
|
||||
|
||||
u32x wordl0[4] = { 0 };
|
||||
u32x wordl1[4] = { 0 };
|
||||
u32x wordl2[4] = { 0 };
|
||||
u32x wordl3[4] = { 0 };
|
||||
|
||||
wordl0[0] = pw_buf0[0];
|
||||
wordl0[1] = pw_buf0[1];
|
||||
wordl0[2] = pw_buf0[2];
|
||||
wordl0[3] = pw_buf0[3];
|
||||
wordl1[0] = pw_buf1[0];
|
||||
wordl1[1] = pw_buf1[1];
|
||||
wordl1[2] = pw_buf1[2];
|
||||
wordl1[3] = pw_buf1[3];
|
||||
|
||||
u32x wordr0[4] = { 0 };
|
||||
u32x wordr1[4] = { 0 };
|
||||
u32x wordr2[4] = { 0 };
|
||||
u32x wordr3[4] = { 0 };
|
||||
|
||||
wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
|
||||
wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
|
||||
wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
|
||||
wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
|
||||
wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
|
||||
wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
|
||||
wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
|
||||
wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
|
||||
|
||||
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
|
||||
{
|
||||
switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
|
||||
}
|
||||
else
|
||||
{
|
||||
switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
|
||||
}
|
||||
|
||||
u32x w0[2];
|
||||
|
||||
w0[0] = wordl0[0] | wordr0[0];
|
||||
w0[1] = wordl0[1] | wordr0[1];
|
||||
|
||||
const u32x c = w0[0];
|
||||
const u32x d = w0[1];
|
||||
|
||||
u32x Kc[16];
|
||||
u32x Kd[16];
|
||||
|
||||
_des_crypt_keysetup (c, d, Kc, Kd, s_skb);
|
||||
|
||||
u32x data[2];
|
||||
|
||||
data[0] = salt_buf0[0];
|
||||
data[1] = salt_buf0[1];
|
||||
|
||||
u32x iv[2];
|
||||
|
||||
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
|
||||
|
||||
u32x z = 0;
|
||||
|
||||
COMPARE_M_SIMD (iv[0], iv[1], z, z);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m14000_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m14000_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m14000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_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);
|
||||
const u32 lid = get_local_id (0);
|
||||
const u32 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* 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;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
u32 pw_buf0[4];
|
||||
u32 pw_buf1[4];
|
||||
|
||||
pw_buf0[0] = pws[gid].i[ 0];
|
||||
pw_buf0[1] = pws[gid].i[ 1];
|
||||
pw_buf0[2] = 0;
|
||||
pw_buf0[3] = 0;
|
||||
pw_buf1[0] = 0;
|
||||
pw_buf1[1] = 0;
|
||||
pw_buf1[2] = 0;
|
||||
pw_buf1[3] = 0;
|
||||
|
||||
const u32 pw_l_len = pws[gid].pw_len;
|
||||
|
||||
/**
|
||||
* salt
|
||||
*/
|
||||
|
||||
u32 salt_buf0[2];
|
||||
|
||||
salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
|
||||
salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
|
||||
|
||||
/**
|
||||
* digest
|
||||
*/
|
||||
|
||||
const u32 search[4] =
|
||||
{
|
||||
digests_buf[digests_offset].digest_buf[DGST_R0],
|
||||
digests_buf[digests_offset].digest_buf[DGST_R1],
|
||||
0,
|
||||
0
|
||||
};
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
|
||||
{
|
||||
const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
|
||||
|
||||
const u32x pw_len = pw_l_len + pw_r_len;
|
||||
|
||||
/**
|
||||
* concat password candidate
|
||||
*/
|
||||
|
||||
u32x wordl0[4] = { 0 };
|
||||
u32x wordl1[4] = { 0 };
|
||||
u32x wordl2[4] = { 0 };
|
||||
u32x wordl3[4] = { 0 };
|
||||
|
||||
wordl0[0] = pw_buf0[0];
|
||||
wordl0[1] = pw_buf0[1];
|
||||
wordl0[2] = pw_buf0[2];
|
||||
wordl0[3] = pw_buf0[3];
|
||||
wordl1[0] = pw_buf1[0];
|
||||
wordl1[1] = pw_buf1[1];
|
||||
wordl1[2] = pw_buf1[2];
|
||||
wordl1[3] = pw_buf1[3];
|
||||
|
||||
u32x wordr0[4] = { 0 };
|
||||
u32x wordr1[4] = { 0 };
|
||||
u32x wordr2[4] = { 0 };
|
||||
u32x wordr3[4] = { 0 };
|
||||
|
||||
wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
|
||||
wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
|
||||
wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
|
||||
wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
|
||||
wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
|
||||
wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
|
||||
wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
|
||||
wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
|
||||
|
||||
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
|
||||
{
|
||||
switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
|
||||
}
|
||||
else
|
||||
{
|
||||
switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
|
||||
}
|
||||
|
||||
u32x w0[2];
|
||||
|
||||
w0[0] = wordl0[0] | wordr0[0];
|
||||
w0[1] = wordl0[1] | wordr0[1];
|
||||
|
||||
|
||||
const u32x c = (w0[0]);
|
||||
const u32x d = (w0[1]);
|
||||
|
||||
u32x Kc[16];
|
||||
u32x Kd[16];
|
||||
|
||||
_des_crypt_keysetup (c, d, Kc, Kd, s_skb);
|
||||
|
||||
u32x data[2];
|
||||
|
||||
data[0] = salt_buf0[0];
|
||||
data[1] = salt_buf0[1];
|
||||
|
||||
u32x iv[2];
|
||||
|
||||
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
|
||||
|
||||
u32x z = 0;
|
||||
|
||||
COMPARE_S_SIMD (iv[0], iv[1], z, z);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m14000_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m14000_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
}
|
||||
2770
OpenCL/m14000_a3.cl
Normal file
2770
OpenCL/m14000_a3.cl
Normal file
File diff suppressed because it is too large
Load Diff
845
OpenCL/m14100_a0.cl
Normal file
845
OpenCL/m14100_a0.cl
Normal file
@ -0,0 +1,845 @@
|
||||
/**
|
||||
* Authors.....: Jens Steube <jens.steube@gmail.com>
|
||||
* Gabriele Gristina <matrix@hashcat.net>
|
||||
* magnum <john.magnum@hushmail.com>
|
||||
* Frans Lategan <frans.lategan+hashcat@gmail.com>
|
||||
*
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
#define _DES_
|
||||
|
||||
#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_rp.h"
|
||||
#include "inc_rp.cl"
|
||||
#include "inc_simd.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 u32 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 u32 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
|
||||
}
|
||||
};
|
||||
|
||||
#if VECT_SIZE == 1
|
||||
#define BOX(i,n,S) (S)[(n)][(i)]
|
||||
#elif VECT_SIZE == 2
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
|
||||
#elif VECT_SIZE == 4
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
|
||||
#elif VECT_SIZE == 8
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7])
|
||||
#elif VECT_SIZE == 16
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
|
||||
#endif
|
||||
|
||||
#if VECT_SIZE == 1
|
||||
#define BOX1(i,S) (S)[(i)]
|
||||
#elif VECT_SIZE == 2
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1])
|
||||
#elif VECT_SIZE == 4
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3])
|
||||
#elif VECT_SIZE == 8
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7])
|
||||
#elif VECT_SIZE == 16
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7], (S)[(i).s8], (S)[(i).s9], (S)[(i).sa], (S)[(i).sb], (S)[(i).sc], (S)[(i).sd], (S)[(i).se], (S)[(i).sf])
|
||||
#endif
|
||||
|
||||
void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
|
||||
{
|
||||
u32x r = rotl32 (data[0], 3u);
|
||||
u32x l = rotl32 (data[1], 3u);
|
||||
|
||||
u32x tt;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i += 2)
|
||||
{
|
||||
u32x u;
|
||||
u32x t;
|
||||
|
||||
u = Kc[i + 0] ^ r;
|
||||
t = Kd[i + 0] ^ rotl32 (r, 28u);
|
||||
|
||||
l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
|
||||
| BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
|
||||
| BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
|
||||
| BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
|
||||
| BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
|
||||
| BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
|
||||
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
|
||||
|
||||
u = Kc[i + 1] ^ l;
|
||||
t = Kd[i + 1] ^ rotl32 (l, 28u);
|
||||
|
||||
r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
|
||||
| BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
|
||||
| BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
|
||||
| BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
|
||||
| BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
|
||||
| BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
|
||||
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
|
||||
}
|
||||
|
||||
iv[0] = rotl32 (l, 29u);
|
||||
iv[1] = rotl32 (r, 29u);
|
||||
}
|
||||
|
||||
void _des_crypt_decrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
|
||||
{
|
||||
u32x r = rotl32 (data[0], 3u);
|
||||
u32x l = rotl32 (data[1], 3u);
|
||||
|
||||
u32x tt;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 16; i > 0; i -= 2)
|
||||
{
|
||||
u32x u;
|
||||
u32x t;
|
||||
|
||||
u = Kc[i - 1] ^ r;
|
||||
t = Kd[i - 1] ^ rotl32 (r, 28u);
|
||||
|
||||
l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
|
||||
| BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
|
||||
| BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
|
||||
| BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
|
||||
| BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
|
||||
| BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
|
||||
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
|
||||
|
||||
u = Kc[i - 2] ^ l;
|
||||
t = Kd[i - 2] ^ rotl32 (l, 28u);
|
||||
|
||||
r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
|
||||
| BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
|
||||
| BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
|
||||
| BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
|
||||
| BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
|
||||
| BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
|
||||
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
|
||||
}
|
||||
|
||||
iv[0] = rotl32 (l, 29u);
|
||||
iv[1] = rotl32 (r, 29u);
|
||||
}
|
||||
|
||||
void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
|
||||
{
|
||||
u32x tt;
|
||||
|
||||
PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
|
||||
HPERM_OP (c, tt, 2, 0xcccc0000);
|
||||
HPERM_OP (d, tt, 2, 0xcccc0000);
|
||||
PERM_OP (d, c, tt, 1, 0x55555555);
|
||||
PERM_OP (c, d, tt, 8, 0x00ff00ff);
|
||||
PERM_OP (d, c, tt, 1, 0x55555555);
|
||||
|
||||
d = ((d & 0x000000ff) << 16)
|
||||
| ((d & 0x0000ff00) << 0)
|
||||
| ((d & 0x00ff0000) >> 16)
|
||||
| ((c & 0xf0000000) >> 4);
|
||||
|
||||
c = c & 0x0fffffff;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i++)
|
||||
{
|
||||
if ((i < 2) || (i == 8) || (i == 15))
|
||||
{
|
||||
c = ((c >> 1) | (c << 27));
|
||||
d = ((d >> 1) | (d << 27));
|
||||
}
|
||||
else
|
||||
{
|
||||
c = ((c >> 2) | (c << 26));
|
||||
d = ((d >> 2) | (d << 26));
|
||||
}
|
||||
|
||||
c = c & 0x0fffffff;
|
||||
d = d & 0x0fffffff;
|
||||
|
||||
const u32x c00 = (c >> 0) & 0x0000003f;
|
||||
const u32x c06 = (c >> 6) & 0x00383003;
|
||||
const u32x c07 = (c >> 7) & 0x0000003c;
|
||||
const u32x c13 = (c >> 13) & 0x0000060f;
|
||||
const u32x c20 = (c >> 20) & 0x00000001;
|
||||
|
||||
u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
|
||||
| BOX (((c06 >> 0) & 0xff)
|
||||
|((c07 >> 0) & 0xff), 1, s_skb)
|
||||
| BOX (((c13 >> 0) & 0xff)
|
||||
|((c06 >> 8) & 0xff), 2, s_skb)
|
||||
| BOX (((c20 >> 0) & 0xff)
|
||||
|((c13 >> 8) & 0xff)
|
||||
|((c06 >> 16) & 0xff), 3, s_skb);
|
||||
|
||||
const u32x d00 = (d >> 0) & 0x00003c3f;
|
||||
const u32x d07 = (d >> 7) & 0x00003f03;
|
||||
const u32x d21 = (d >> 21) & 0x0000000f;
|
||||
const u32x d22 = (d >> 22) & 0x00000030;
|
||||
|
||||
u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
|
||||
| BOX (((d07 >> 0) & 0xff)
|
||||
|((d00 >> 8) & 0xff), 5, s_skb)
|
||||
| BOX (((d07 >> 8) & 0xff), 6, s_skb)
|
||||
| BOX (((d21 >> 0) & 0xff)
|
||||
|((d22 >> 0) & 0xff), 7, s_skb);
|
||||
|
||||
Kc[i] = ((t << 16) | (s & 0x0000ffff));
|
||||
Kd[i] = ((s >> 16) | (t & 0xffff0000));
|
||||
|
||||
Kc[i] = rotl32 (Kc[i], 2u);
|
||||
Kd[i] = rotl32 (Kd[i], 2u);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m14100_m04 (__global pw_t *pws, __global kernel_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_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);
|
||||
const u32 lid = get_local_id (0);
|
||||
const u32 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* 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;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
u32 pw_buf0[4];
|
||||
u32 pw_buf1[4];
|
||||
|
||||
pw_buf0[0] = pws[gid].i[ 0];
|
||||
pw_buf0[1] = pws[gid].i[ 1];
|
||||
pw_buf0[2] = pws[gid].i[ 2];
|
||||
pw_buf0[3] = pws[gid].i[ 3];
|
||||
pw_buf1[0] = pws[gid].i[ 4];
|
||||
pw_buf1[1] = pws[gid].i[ 5];
|
||||
pw_buf1[2] = 0;
|
||||
pw_buf1[3] = 0;
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
/**
|
||||
* salt
|
||||
*/
|
||||
|
||||
u32 salt_buf0[2];
|
||||
|
||||
salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
|
||||
salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
|
||||
|
||||
/**
|
||||
* main
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
|
||||
{
|
||||
u32x w0[4] = { 0 };
|
||||
u32x w1[4] = { 0 };
|
||||
u32x w2[4] = { 0 };
|
||||
u32x w3[4] = { 0 };
|
||||
|
||||
apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
|
||||
|
||||
/* First Pass */
|
||||
|
||||
const u32x a = w0[0];
|
||||
const u32x b = w0[1];
|
||||
|
||||
u32x Ka[16];
|
||||
u32x Kb[16];
|
||||
|
||||
_des_crypt_keysetup (a, b, Ka, Kb, s_skb);
|
||||
|
||||
u32x data[2];
|
||||
|
||||
data[0] = salt_buf0[0];
|
||||
data[1] = salt_buf0[1];
|
||||
|
||||
u32x p1[2];
|
||||
|
||||
_des_crypt_encrypt (p1, data, Ka, Kb, s_SPtrans);
|
||||
|
||||
/* Second Pass */
|
||||
|
||||
const u32x c = w0[2];
|
||||
const u32x d = w0[3];
|
||||
|
||||
u32x Kc[16];
|
||||
u32x Kd[16];
|
||||
|
||||
_des_crypt_keysetup (c, d, Kc, Kd, s_skb);
|
||||
|
||||
u32x p2[2];
|
||||
|
||||
_des_crypt_decrypt (p2, p1, Kc, Kd, s_SPtrans);
|
||||
|
||||
/* Third Pass */
|
||||
|
||||
const u32x e = w1[0];
|
||||
const u32x f = w1[1];
|
||||
|
||||
u32x Ke[16];
|
||||
u32x Kf[16];
|
||||
|
||||
_des_crypt_keysetup (e, f, Ke, Kf, s_skb);
|
||||
|
||||
u32x iv[2];
|
||||
|
||||
_des_crypt_encrypt (iv, p2, Ke, Kf, s_SPtrans);
|
||||
|
||||
u32x z = 0;
|
||||
|
||||
COMPARE_M_SIMD (iv[0], iv[1], z, z);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m14100_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m14100_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m14100_s04 (__global pw_t *pws, __global kernel_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_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);
|
||||
const u32 lid = get_local_id (0);
|
||||
const u32 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* 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;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
u32 pw_buf0[4];
|
||||
u32 pw_buf1[4];
|
||||
|
||||
pw_buf0[0] = pws[gid].i[ 0];
|
||||
pw_buf0[1] = pws[gid].i[ 1];
|
||||
pw_buf0[2] = pws[gid].i[ 2];
|
||||
pw_buf0[3] = pws[gid].i[ 3];
|
||||
pw_buf1[0] = pws[gid].i[ 4];
|
||||
pw_buf1[1] = pws[gid].i[ 5];
|
||||
pw_buf1[2] = 0;
|
||||
pw_buf1[3] = 0;
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
/**
|
||||
* salt
|
||||
*/
|
||||
|
||||
u32 salt_buf0[2];
|
||||
|
||||
salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
|
||||
salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
|
||||
|
||||
/**
|
||||
* digest
|
||||
*/
|
||||
|
||||
const u32 search[4] =
|
||||
{
|
||||
digests_buf[digests_offset].digest_buf[DGST_R0],
|
||||
digests_buf[digests_offset].digest_buf[DGST_R1],
|
||||
0,
|
||||
0
|
||||
};
|
||||
|
||||
/**
|
||||
* main
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
|
||||
{
|
||||
u32x w0[4] = { 0 };
|
||||
u32x w1[4] = { 0 };
|
||||
u32x w2[4] = { 0 };
|
||||
u32x w3[4] = { 0 };
|
||||
|
||||
apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
|
||||
|
||||
/* First Pass */
|
||||
|
||||
const u32x a = w0[0];
|
||||
const u32x b = w0[1];
|
||||
|
||||
u32x Ka[16];
|
||||
u32x Kb[16];
|
||||
|
||||
_des_crypt_keysetup (a, b, Ka, Kb, s_skb);
|
||||
|
||||
u32x data[2];
|
||||
|
||||
data[0] = salt_buf0[0];
|
||||
data[1] = salt_buf0[1];
|
||||
|
||||
u32x p1[2];
|
||||
|
||||
_des_crypt_encrypt (p1, data, Ka, Kb, s_SPtrans);
|
||||
|
||||
/* Second Pass */
|
||||
|
||||
const u32x c = w0[2];
|
||||
const u32x d = w0[3];
|
||||
|
||||
u32x Kc[16];
|
||||
u32x Kd[16];
|
||||
|
||||
_des_crypt_keysetup (c, d, Kc, Kd, s_skb);
|
||||
|
||||
u32x p2[2];
|
||||
|
||||
_des_crypt_decrypt (p2, p1, Kc, Kd, s_SPtrans);
|
||||
|
||||
/* Third Pass */
|
||||
|
||||
const u32x e = w1[0];
|
||||
const u32x f = w1[1];
|
||||
|
||||
u32x Ke[16];
|
||||
u32x Kf[16];
|
||||
|
||||
_des_crypt_keysetup (e, f, Ke, Kf, s_skb);
|
||||
|
||||
u32x iv[2];
|
||||
|
||||
_des_crypt_encrypt (iv, p2, Ke, Kf, s_SPtrans);
|
||||
|
||||
u32x z = 0;
|
||||
|
||||
COMPARE_S_SIMD (iv[0], iv[1], z, z);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m14100_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m14100_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
}
|
||||
932
OpenCL/m14100_a1.cl
Normal file
932
OpenCL/m14100_a1.cl
Normal file
@ -0,0 +1,932 @@
|
||||
/**
|
||||
* Authors.....: Jens Steube <jens.steube@gmail.com>
|
||||
* Gabriele Gristina <matrix@hashcat.net>
|
||||
* Frans Lategan <frans.lategan+hashcat@gmail.com>
|
||||
*
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
#define _DES_
|
||||
|
||||
#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"
|
||||
|
||||
#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 u32 c_SPtrans[8][64] =
|
||||
{
|
||||
{
|
||||
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,
|
||||
},
|
||||
{
|
||||
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,
|
||||
},
|
||||
{
|
||||
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,
|
||||
},
|
||||
{
|
||||
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,
|
||||
},
|
||||
{
|
||||
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,
|
||||
},
|
||||
{
|
||||
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,
|
||||
},
|
||||
{
|
||||
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,
|
||||
},
|
||||
{
|
||||
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 u32 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
|
||||
}
|
||||
};
|
||||
|
||||
#if VECT_SIZE == 1
|
||||
#define BOX(i,n,S) (S)[(n)][(i)]
|
||||
#elif VECT_SIZE == 2
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
|
||||
#elif VECT_SIZE == 4
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
|
||||
#elif VECT_SIZE == 8
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7])
|
||||
#elif VECT_SIZE == 16
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
|
||||
#endif
|
||||
|
||||
#if VECT_SIZE == 1
|
||||
#define BOX1(i,S) (S)[(i)]
|
||||
#elif VECT_SIZE == 2
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1])
|
||||
#elif VECT_SIZE == 4
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3])
|
||||
#elif VECT_SIZE == 8
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7])
|
||||
#elif VECT_SIZE == 16
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7], (S)[(i).s8], (S)[(i).s9], (S)[(i).sa], (S)[(i).sb], (S)[(i).sc], (S)[(i).sd], (S)[(i).se], (S)[(i).sf])
|
||||
#endif
|
||||
|
||||
void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
|
||||
{
|
||||
u32x r = rotl32 (data[0], 3u);
|
||||
u32x l = rotl32 (data[1], 3u);
|
||||
|
||||
u32x tt;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i += 2)
|
||||
{
|
||||
u32x u;
|
||||
u32x t;
|
||||
|
||||
u = Kc[i + 0] ^ r;
|
||||
t = Kd[i + 0] ^ rotl32 (r, 28u);
|
||||
|
||||
l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
|
||||
| BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
|
||||
| BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
|
||||
| BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
|
||||
| BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
|
||||
| BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
|
||||
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
|
||||
|
||||
u = Kc[i + 1] ^ l;
|
||||
t = Kd[i + 1] ^ rotl32 (l, 28u);
|
||||
|
||||
r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
|
||||
| BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
|
||||
| BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
|
||||
| BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
|
||||
| BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
|
||||
| BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
|
||||
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
|
||||
}
|
||||
|
||||
iv[0] = rotl32 (l, 29u);
|
||||
iv[1] = rotl32 (r, 29u);
|
||||
}
|
||||
|
||||
void _des_crypt_decrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
|
||||
{
|
||||
u32x r = rotl32 (data[0], 3u);
|
||||
u32x l = rotl32 (data[1], 3u);
|
||||
|
||||
u32x tt;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 16; i > 0; i -= 2)
|
||||
{
|
||||
u32x u;
|
||||
u32x t;
|
||||
|
||||
u = Kc[i - 1] ^ r;
|
||||
t = Kd[i - 1] ^ rotl32 (r, 28u);
|
||||
|
||||
l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
|
||||
| BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
|
||||
| BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
|
||||
| BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
|
||||
| BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
|
||||
| BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
|
||||
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
|
||||
|
||||
u = Kc[i - 2] ^ l;
|
||||
t = Kd[i - 2] ^ rotl32 (l, 28u);
|
||||
|
||||
r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
|
||||
| BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
|
||||
| BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
|
||||
| BOX (((t >> 2) & 0x3f), 1, s_SPtrans)
|
||||
| BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
|
||||
| BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
|
||||
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
|
||||
}
|
||||
|
||||
iv[0] = rotl32 (l, 29u);
|
||||
iv[1] = rotl32 (r, 29u);
|
||||
}
|
||||
|
||||
void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
|
||||
{
|
||||
u32x tt;
|
||||
|
||||
PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
|
||||
HPERM_OP (c, tt, 2, 0xcccc0000);
|
||||
HPERM_OP (d, tt, 2, 0xcccc0000);
|
||||
PERM_OP (d, c, tt, 1, 0x55555555);
|
||||
PERM_OP (c, d, tt, 8, 0x00ff00ff);
|
||||
PERM_OP (d, c, tt, 1, 0x55555555);
|
||||
|
||||
d = ((d & 0x000000ff) << 16)
|
||||
| ((d & 0x0000ff00) << 0)
|
||||
| ((d & 0x00ff0000) >> 16)
|
||||
| ((c & 0xf0000000) >> 4);
|
||||
|
||||
c = c & 0x0fffffff;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i++)
|
||||
{
|
||||
if ((i < 2) || (i == 8) || (i == 15))
|
||||
{
|
||||
c = ((c >> 1) | (c << 27));
|
||||
d = ((d >> 1) | (d << 27));
|
||||
}
|
||||
else
|
||||
{
|
||||
c = ((c >> 2) | (c << 26));
|
||||
d = ((d >> 2) | (d << 26));
|
||||
}
|
||||
|
||||
c = c & 0x0fffffff;
|
||||
d = d & 0x0fffffff;
|
||||
|
||||
const u32x c00 = (c >> 0) & 0x0000003f;
|
||||
const u32x c06 = (c >> 6) & 0x00383003;
|
||||
const u32x c07 = (c >> 7) & 0x0000003c;
|
||||
const u32x c13 = (c >> 13) & 0x0000060f;
|
||||
const u32x c20 = (c >> 20) & 0x00000001;
|
||||
|
||||
u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
|
||||
| BOX (((c06 >> 0) & 0xff)
|
||||
|((c07 >> 0) & 0xff), 1, s_skb)
|
||||
| BOX (((c13 >> 0) & 0xff)
|
||||
|((c06 >> 8) & 0xff), 2, s_skb)
|
||||
| BOX (((c20 >> 0) & 0xff)
|
||||
|((c13 >> 8) & 0xff)
|
||||
|((c06 >> 16) & 0xff), 3, s_skb);
|
||||
|
||||
const u32x d00 = (d >> 0) & 0x00003c3f;
|
||||
const u32x d07 = (d >> 7) & 0x00003f03;
|
||||
const u32x d21 = (d >> 21) & 0x0000000f;
|
||||
const u32x d22 = (d >> 22) & 0x00000030;
|
||||
|
||||
u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
|
||||
| BOX (((d07 >> 0) & 0xff)
|
||||
|((d00 >> 8) & 0xff), 5, s_skb)
|
||||
| BOX (((d07 >> 8) & 0xff), 6, s_skb)
|
||||
| BOX (((d21 >> 0) & 0xff)
|
||||
|((d22 >> 0) & 0xff), 7, s_skb);
|
||||
|
||||
Kc[i] = ((t << 16) | (s & 0x0000ffff));
|
||||
Kd[i] = ((s >> 16) | (t & 0xffff0000));
|
||||
|
||||
Kc[i] = rotl32 (Kc[i], 2u);
|
||||
Kd[i] = rotl32 (Kd[i], 2u);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m14100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_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);
|
||||
const u32 lid = get_local_id (0);
|
||||
const u32 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* 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;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
u32 pw_buf0[4];
|
||||
u32 pw_buf1[4];
|
||||
|
||||
pw_buf0[0] = pws[gid].i[ 0];
|
||||
pw_buf0[1] = pws[gid].i[ 1];
|
||||
pw_buf0[2] = pws[gid].i[ 2];
|
||||
pw_buf0[3] = pws[gid].i[ 3];
|
||||
pw_buf1[0] = pws[gid].i[ 4];
|
||||
pw_buf1[1] = pws[gid].i[ 5];
|
||||
pw_buf1[2] = 0;
|
||||
pw_buf1[3] = 0;
|
||||
|
||||
|
||||
const u32 pw_l_len = pws[gid].pw_len;
|
||||
|
||||
/**
|
||||
* salt
|
||||
*/
|
||||
|
||||
u32 salt_buf0[2];
|
||||
|
||||
salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
|
||||
salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
|
||||
{
|
||||
const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
|
||||
|
||||
const u32x pw_len = pw_l_len + pw_r_len;
|
||||
|
||||
/**
|
||||
* concat password candidate
|
||||
*/
|
||||
|
||||
u32x wordl0[4] = { 0 };
|
||||
u32x wordl1[4] = { 0 };
|
||||
u32x wordl2[4] = { 0 };
|
||||
u32x wordl3[4] = { 0 };
|
||||
|
||||
wordl0[0] = pw_buf0[0];
|
||||
wordl0[1] = pw_buf0[1];
|
||||
wordl0[2] = pw_buf0[2];
|
||||
wordl0[3] = pw_buf0[3];
|
||||
wordl1[0] = pw_buf1[0];
|
||||
wordl1[1] = pw_buf1[1];
|
||||
wordl1[2] = pw_buf1[2];
|
||||
wordl1[3] = pw_buf1[3];
|
||||
|
||||
u32x wordr0[4] = { 0 };
|
||||
u32x wordr1[4] = { 0 };
|
||||
u32x wordr2[4] = { 0 };
|
||||
u32x wordr3[4] = { 0 };
|
||||
|
||||
wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
|
||||
wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
|
||||
wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
|
||||
wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
|
||||
wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
|
||||
wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
|
||||
wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
|
||||
wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
|
||||
|
||||
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
|
||||
{
|
||||
switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
|
||||
}
|
||||
else
|
||||
{
|
||||
switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
|
||||
}
|
||||
|
||||
u32x w0[4];
|
||||
u32x w1[4];
|
||||
|
||||
w0[0] = wordl0[0] | wordr0[0];
|
||||
w0[1] = wordl0[1] | wordr0[1];
|
||||
w0[2] = wordl0[2] | wordr0[2];
|
||||
w0[3] = wordl0[3] | wordr0[3];
|
||||
w1[0] = wordl1[0] | wordr1[0];
|
||||
w1[1] = wordl1[1] | wordr1[1];
|
||||
|
||||
/* First Pass */
|
||||
|
||||
const u32x a = w0[0];
|
||||
const u32x b = w0[1];
|
||||
|
||||
u32x Ka[16];
|
||||
u32x Kb[16];
|
||||
|
||||
_des_crypt_keysetup (a, b, Ka, Kb, s_skb);
|
||||
|
||||
u32x data[2];
|
||||
|
||||
data[0] = salt_buf0[0];
|
||||
data[1] = salt_buf0[1];
|
||||
|
||||
u32x p1[2];
|
||||
|
||||
_des_crypt_encrypt (p1, data, Ka, Kb, s_SPtrans);
|
||||
|
||||
/* Second Pass */
|
||||
|
||||
const u32x c = w0[2];
|
||||
const u32x d = w0[3];
|
||||
|
||||
u32x Kc[16];
|
||||
u32x Kd[16];
|
||||
|
||||
_des_crypt_keysetup (c, d, Kc, Kd, s_skb);
|
||||
|
||||
u32x p2[2];
|
||||
|
||||
_des_crypt_decrypt (p2, p1, Kc, Kd, s_SPtrans);
|
||||
|
||||
/* Third Pass */
|
||||
|
||||
const u32x e = w1[0];
|
||||
const u32x f = w1[1];
|
||||
|
||||
u32x Ke[16];
|
||||
u32x Kf[16];
|
||||
|
||||
_des_crypt_keysetup (e, f, Ke, Kf, s_skb);
|
||||
|
||||
u32x iv[2];
|
||||
|
||||
_des_crypt_encrypt (iv, p2, Ke, Kf, s_SPtrans);
|
||||
|
||||
u32x z = 0;
|
||||
|
||||
COMPARE_M_SIMD (iv[0], iv[1], z, z);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m14100_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m14100_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m14100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_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);
|
||||
const u32 lid = get_local_id (0);
|
||||
const u32 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* 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;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
u32 pw_buf0[4];
|
||||
u32 pw_buf1[4];
|
||||
|
||||
pw_buf0[0] = pws[gid].i[ 0];
|
||||
pw_buf0[1] = pws[gid].i[ 1];
|
||||
pw_buf0[2] = pws[gid].i[ 2];
|
||||
pw_buf0[3] = pws[gid].i[ 3];
|
||||
pw_buf1[0] = pws[gid].i[ 4];
|
||||
pw_buf1[1] = pws[gid].i[ 5];
|
||||
pw_buf1[2] = 0;
|
||||
pw_buf1[3] = 0;
|
||||
|
||||
|
||||
const u32 pw_l_len = pws[gid].pw_len;
|
||||
|
||||
/**
|
||||
* salt
|
||||
*/
|
||||
|
||||
u32 salt_buf0[2];
|
||||
|
||||
salt_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[0];
|
||||
salt_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[1];
|
||||
|
||||
/**
|
||||
* digest
|
||||
*/
|
||||
|
||||
const u32 search[4] =
|
||||
{
|
||||
digests_buf[digests_offset].digest_buf[DGST_R0],
|
||||
digests_buf[digests_offset].digest_buf[DGST_R1],
|
||||
0,
|
||||
0
|
||||
};
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
|
||||
{
|
||||
const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
|
||||
|
||||
const u32x pw_len = pw_l_len + pw_r_len;
|
||||
|
||||
/**
|
||||
* concat password candidate
|
||||
*/
|
||||
|
||||
u32x wordl0[4] = { 0 };
|
||||
u32x wordl1[4] = { 0 };
|
||||
u32x wordl2[4] = { 0 };
|
||||
u32x wordl3[4] = { 0 };
|
||||
|
||||
wordl0[0] = pw_buf0[0];
|
||||
wordl0[1] = pw_buf0[1];
|
||||
wordl0[2] = pw_buf0[2];
|
||||
wordl0[3] = pw_buf0[3];
|
||||
wordl1[0] = pw_buf1[0];
|
||||
wordl1[1] = pw_buf1[1];
|
||||
wordl1[2] = pw_buf1[2];
|
||||
wordl1[3] = pw_buf1[3];
|
||||
|
||||
u32x wordr0[4] = { 0 };
|
||||
u32x wordr1[4] = { 0 };
|
||||
u32x wordr2[4] = { 0 };
|
||||
u32x wordr3[4] = { 0 };
|
||||
|
||||
wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
|
||||
wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
|
||||
wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
|
||||
wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
|
||||
wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
|
||||
wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
|
||||
wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
|
||||
wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
|
||||
|
||||
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
|
||||
{
|
||||
switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
|
||||
}
|
||||
else
|
||||
{
|
||||
switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
|
||||
}
|
||||
|
||||
u32x w0[4];
|
||||
u32x w1[4];
|
||||
|
||||
w0[0] = wordl0[0] | wordr0[0];
|
||||
w0[1] = wordl0[1] | wordr0[1];
|
||||
w0[2] = wordl0[2] | wordr0[2];
|
||||
w0[3] = wordl0[3] | wordr0[3];
|
||||
w1[0] = wordl1[0] | wordr1[0];
|
||||
w1[1] = wordl1[1] | wordr1[1];
|
||||
|
||||
/* First Pass */
|
||||
|
||||
const u32x a = w0[0];
|
||||
const u32x b = w0[1];
|
||||
|
||||
u32x Ka[16];
|
||||
u32x Kb[16];
|
||||
|
||||
_des_crypt_keysetup (a, b, Ka, Kb, s_skb);
|
||||
|
||||
u32x data[2];
|
||||
|
||||
data[0] = salt_buf0[0];
|
||||
data[1] = salt_buf0[1];
|
||||
|
||||
u32x p1[2];
|
||||
|
||||
_des_crypt_encrypt (p1, data, Ka, Kb, s_SPtrans);
|
||||
|
||||
/* Second Pass */
|
||||
|
||||
const u32x c = w0[2];
|
||||
const u32x d = w0[3];
|
||||
|
||||
u32x Kc[16];
|
||||
u32x Kd[16];
|
||||
|
||||
_des_crypt_keysetup (c, d, Kc, Kd, s_skb);
|
||||
|
||||
u32x p2[2];
|
||||
|
||||
_des_crypt_decrypt (p2, p1, Kc, Kd, s_SPtrans);
|
||||
|
||||
/* Third Pass */
|
||||
|
||||
const u32x e = w1[0];
|
||||
const u32x f = w1[1];
|
||||
|
||||
u32x Ke[16];
|
||||
u32x Kf[16];
|
||||
|
||||
_des_crypt_keysetup (e, f, Ke, Kf, s_skb);
|
||||
|
||||
u32x iv[2];
|
||||
|
||||
_des_crypt_encrypt (iv, p2, Ke, Kf, s_SPtrans);
|
||||
|
||||
u32x z = 0;
|
||||
|
||||
COMPARE_S_SIMD (iv[0], iv[1], z, z);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m14100_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m14100_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
}
|
||||
1163
OpenCL/m14100_a3.cl
Normal file
1163
OpenCL/m14100_a3.cl
Normal file
File diff suppressed because it is too large
Load Diff
@ -1,3 +1,12 @@
|
||||
* changes v3.00 -> v3.xx:
|
||||
|
||||
##
|
||||
## Algorithms
|
||||
##
|
||||
|
||||
- Added new hash-mode 14000 = DES (PT = $salt, key = $pass)
|
||||
- Added new hash-mode 14100 = 3DES (PT = $salt, key = $pass)
|
||||
|
||||
* changes v3.00 -> v3.10:
|
||||
|
||||
##
|
||||
|
||||
@ -45,6 +45,8 @@ NVidia users require NVidia drivers 346.59 or later (recommended 367.27 or later
|
||||
- SipHash
|
||||
- RipeMD160
|
||||
- Whirlpool
|
||||
- DES (PT = $salt, key = $pass)
|
||||
- 3DES (PT = $salt, key = $pass)
|
||||
- GOST R 34.11-94
|
||||
- GOST R 34.11-2012 (Streebog) 256-bit
|
||||
- GOST R 34.11-2012 (Streebog) 512-bit
|
||||
|
||||
@ -370,6 +370,8 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define HT_13600 "WinZip"
|
||||
#define HT_13800 "Windows 8+ phone PIN/Password"
|
||||
#define HT_13900 "OpenCart"
|
||||
#define HT_14000 "DES (PT = $salt, key = $pass)"
|
||||
#define HT_14100 "3DES (PT = $salt, key = $pass)"
|
||||
|
||||
#define HT_00011 "Joomla < 2.5.18"
|
||||
#define HT_00012 "PostgreSQL"
|
||||
@ -736,6 +738,10 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define DISPLAY_LEN_MAX_13800 64 + 1 + 256
|
||||
#define DISPLAY_LEN_MIN_13900 40 + 1 + 9
|
||||
#define DISPLAY_LEN_MAX_13900 40 + 1 + 9
|
||||
#define DISPLAY_LEN_MIN_14000 16 + 1 + 16
|
||||
#define DISPLAY_LEN_MAX_14000 16 + 1 + 16
|
||||
#define DISPLAY_LEN_MIN_14100 16 + 1 + 16
|
||||
#define DISPLAY_LEN_MAX_14100 16 + 1 + 16
|
||||
|
||||
#define DISPLAY_LEN_MIN_11 32 + 1 + 16
|
||||
#define DISPLAY_LEN_MAX_11 32 + 1 + 32
|
||||
@ -853,6 +859,7 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define HASH_TYPE_RAR3HP 49
|
||||
#define HASH_TYPE_KRB5TGS 50
|
||||
#define HASH_TYPE_STDOUT 51
|
||||
#define HASH_TYPE_DES 52
|
||||
|
||||
#define KERN_TYPE_MD5 0
|
||||
#define KERN_TYPE_MD5_PWSLT 10
|
||||
@ -1009,6 +1016,8 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define KERN_TYPE_ZIP2 13600
|
||||
#define KERN_TYPE_WIN8PHONE 13800
|
||||
#define KERN_TYPE_OPENCART 13900
|
||||
#define KERN_TYPE_DES 14000
|
||||
#define KERN_TYPE_3DES 14100
|
||||
|
||||
/**
|
||||
* signatures
|
||||
@ -1522,6 +1531,7 @@ int cisco4_parse_hash (char *input_buf, uint input_len, hash_t *hash
|
||||
int dcc_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf);
|
||||
int dcc2_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf);
|
||||
int descrypt_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf);
|
||||
int des_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf);
|
||||
int episerver_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf);
|
||||
int ipb2_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf);
|
||||
int joomla_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf);
|
||||
|
||||
@ -154,7 +154,7 @@ double TARGET_MS_PROFILE[4] = { 2, 12, 96, 480 };
|
||||
|
||||
#define MAX_DICTSTAT 10000
|
||||
|
||||
#define NUM_DEFAULT_BENCHMARK_ALGORITHMS 144
|
||||
#define NUM_DEFAULT_BENCHMARK_ALGORITHMS 146
|
||||
|
||||
#define NVIDIA_100PERCENTCPU_WORKAROUND 100
|
||||
|
||||
@ -194,6 +194,8 @@ static uint default_benchmark_algorithms[NUM_DEFAULT_BENCHMARK_ALGORITHMS] =
|
||||
6900,
|
||||
11700,
|
||||
11800,
|
||||
14000,
|
||||
14100,
|
||||
400,
|
||||
8900,
|
||||
11900,
|
||||
@ -499,6 +501,8 @@ const char *USAGE_BIG[] =
|
||||
" 1460 | HMAC-SHA256 (key = $salt) | Raw Hash, Authenticated",
|
||||
" 1750 | HMAC-SHA512 (key = $pass) | Raw Hash, Authenticated",
|
||||
" 1760 | HMAC-SHA512 (key = $salt) | Raw Hash, Authenticated",
|
||||
" 14000 | DES (PT = $salt, key = $pass) | Raw Cipher, Known-Plaintext attack",
|
||||
" 14100 | 3DES (PT = $salt, key = $pass) | Raw Cipher, Known-Plaintext attack",
|
||||
" 400 | phpass | Generic KDF",
|
||||
" 8900 | scrypt | Generic KDF",
|
||||
" 11900 | PBKDF2-HMAC-MD5 | Generic KDF",
|
||||
@ -6999,7 +7003,7 @@ int main (int argc, char **argv)
|
||||
return -1;
|
||||
}
|
||||
|
||||
if (hash_mode_chgd && hash_mode > 13900) // just added to remove compiler warnings for hash_mode_chgd
|
||||
if (hash_mode_chgd && hash_mode > 14100) // just added to remove compiler warnings for hash_mode_chgd
|
||||
{
|
||||
log_error ("ERROR: Invalid hash-type specified");
|
||||
|
||||
@ -11794,6 +11798,43 @@ int main (int argc, char **argv)
|
||||
dgst_pos3 = 1;
|
||||
break;
|
||||
|
||||
case 14000: hash_type = HASH_TYPE_DES;
|
||||
salt_type = SALT_TYPE_EMBEDDED;
|
||||
attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
|
||||
opts_type = OPTS_TYPE_PT_GENERATE_LE
|
||||
| OPTS_TYPE_PT_BITSLICE
|
||||
| OPTS_TYPE_ST_GENERATE_LE
|
||||
| OPTS_TYPE_ST_HEX;
|
||||
kern_type = KERN_TYPE_DES;
|
||||
dgst_size = DGST_SIZE_4_4; // originally DGST_SIZE_4_2
|
||||
parse_func = des_parse_hash;
|
||||
sort_by_digest = sort_by_digest_4_4; // originally sort_by_digest_4_2
|
||||
opti_type = OPTI_TYPE_ZERO_BYTE
|
||||
| OPTI_TYPE_PRECOMPUTE_PERMUT;
|
||||
dgst_pos0 = 0;
|
||||
dgst_pos1 = 1;
|
||||
dgst_pos2 = 2;
|
||||
dgst_pos3 = 3;
|
||||
break;
|
||||
|
||||
case 14100: hash_type = HASH_TYPE_DES;
|
||||
salt_type = SALT_TYPE_EMBEDDED;
|
||||
attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
|
||||
opts_type = OPTS_TYPE_PT_GENERATE_LE
|
||||
| OPTS_TYPE_ST_GENERATE_LE
|
||||
| OPTS_TYPE_ST_HEX;
|
||||
kern_type = KERN_TYPE_3DES;
|
||||
dgst_size = DGST_SIZE_4_4; // originally DGST_SIZE_4_2
|
||||
parse_func = des_parse_hash;
|
||||
sort_by_digest = sort_by_digest_4_4; // originally sort_by_digest_4_2
|
||||
opti_type = OPTI_TYPE_ZERO_BYTE
|
||||
| OPTI_TYPE_PRECOMPUTE_PERMUT;
|
||||
dgst_pos0 = 0;
|
||||
dgst_pos1 = 1;
|
||||
dgst_pos2 = 2;
|
||||
dgst_pos3 = 3;
|
||||
break;
|
||||
|
||||
default: usage_mini_print (PROGNAME); return -1;
|
||||
}
|
||||
|
||||
@ -12286,6 +12327,12 @@ int main (int argc, char **argv)
|
||||
break;
|
||||
case 12800: if (pw_max > 24) pw_max = 24;
|
||||
break;
|
||||
case 14000: if (pw_min < 8) pw_min = 8;
|
||||
if (pw_max > 8) pw_max = 8;
|
||||
break;
|
||||
case 14100: if (pw_min < 24) pw_min = 24;
|
||||
if (pw_max > 24) pw_max = 24;
|
||||
break;
|
||||
}
|
||||
|
||||
if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
|
||||
@ -12978,6 +13025,11 @@ int main (int argc, char **argv)
|
||||
break;
|
||||
case 12600: hashes_buf[0].salt->salt_len = 64;
|
||||
break;
|
||||
case 14000: hashes_buf[0].salt->salt_len = 8;
|
||||
break;
|
||||
case 14100: hashes_buf[0].salt->salt_len = 8;
|
||||
break;
|
||||
|
||||
}
|
||||
|
||||
// special esalt handling
|
||||
@ -15871,8 +15923,10 @@ int main (int argc, char **argv)
|
||||
|
||||
if (hash_mode == 1500) kernel_threads = 64; // DES
|
||||
if (hash_mode == 3000) kernel_threads = 64; // DES
|
||||
if (hash_mode == 3100) kernel_threads = 64; // DES
|
||||
if (hash_mode == 3200) kernel_threads = 8; // Blowfish
|
||||
if (hash_mode == 7500) kernel_threads = 64; // RC4
|
||||
if (hash_mode == 8500) kernel_threads = 64; // DES
|
||||
if (hash_mode == 9000) kernel_threads = 8; // Blowfish
|
||||
if (hash_mode == 9700) kernel_threads = 64; // RC4
|
||||
if (hash_mode == 9710) kernel_threads = 64; // RC4
|
||||
@ -15882,6 +15936,8 @@ int main (int argc, char **argv)
|
||||
if (hash_mode == 10410) kernel_threads = 64; // RC4
|
||||
if (hash_mode == 10500) kernel_threads = 64; // RC4
|
||||
if (hash_mode == 13100) kernel_threads = 64; // RC4
|
||||
if (hash_mode == 14000) kernel_threads = 64; // DES
|
||||
if (hash_mode == 14100) kernel_threads = 64; // DES
|
||||
|
||||
device_param->kernel_threads = kernel_threads;
|
||||
|
||||
@ -16070,6 +16126,22 @@ int main (int argc, char **argv)
|
||||
device_param->kernel_loops_max = kernel_loops_fixed;
|
||||
}
|
||||
|
||||
if (hash_mode == 14000 && attack_mode == ATTACK_MODE_BF)
|
||||
{
|
||||
const u32 kernel_loops_fixed = 1024;
|
||||
|
||||
device_param->kernel_loops_min = kernel_loops_fixed;
|
||||
device_param->kernel_loops_max = kernel_loops_fixed;
|
||||
}
|
||||
|
||||
if (hash_mode == 14100 && attack_mode == ATTACK_MODE_BF)
|
||||
{
|
||||
const u32 kernel_loops_fixed = 1024;
|
||||
|
||||
device_param->kernel_loops_min = kernel_loops_fixed;
|
||||
device_param->kernel_loops_max = kernel_loops_fixed;
|
||||
}
|
||||
|
||||
/**
|
||||
* some algorithms have a maximum kernel-loops count
|
||||
*/
|
||||
|
||||
69
src/shared.c
69
src/shared.c
@ -6148,6 +6148,8 @@ char *strhashtype (const uint hash_mode)
|
||||
case 13763: return ((char *) HT_13763);
|
||||
case 13800: return ((char *) HT_13800);
|
||||
case 13900: return ((char *) HT_13900);
|
||||
case 14000: return ((char *) HT_14000);
|
||||
case 14100: return ((char *) HT_14100);
|
||||
}
|
||||
|
||||
return ((char *) "Unknown");
|
||||
@ -6201,6 +6203,10 @@ void ascii_digest (char *out_buf, uint salt_pos, uint digest_pos)
|
||||
|
||||
switch (hash_type)
|
||||
{
|
||||
case HASH_TYPE_DES:
|
||||
FP (digest_buf[1], digest_buf[0], tt);
|
||||
break;
|
||||
|
||||
case HASH_TYPE_DESCRYPT:
|
||||
FP (digest_buf[1], digest_buf[0], tt);
|
||||
break;
|
||||
@ -8879,6 +8885,14 @@ void ascii_digest (char *out_buf, uint salt_pos, uint digest_pos)
|
||||
digest_buf[7],
|
||||
buf);
|
||||
}
|
||||
else if (hash_mode == 14000)
|
||||
{
|
||||
snprintf (out_buf, len - 1, "%08x%08x:%s", digest_buf[0], digest_buf[1], (char *) salt.salt_buf);
|
||||
}
|
||||
else if (hash_mode == 14100)
|
||||
{
|
||||
snprintf (out_buf, len - 1, "%08x%08x:%s", digest_buf[0], digest_buf[1], (char *) salt.salt_buf);
|
||||
}
|
||||
else
|
||||
{
|
||||
if (hash_type == HASH_TYPE_MD4)
|
||||
@ -15098,6 +15112,61 @@ int racf_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf)
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
int des_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf)
|
||||
{
|
||||
if ((input_len < DISPLAY_LEN_MIN_14000) || (input_len > DISPLAY_LEN_MAX_14000)) return (PARSER_GLOBAL_LENGTH);
|
||||
|
||||
u32 *digest = (u32 *) hash_buf->digest;
|
||||
|
||||
salt_t *salt = hash_buf->salt;
|
||||
|
||||
char *digest_pos = input_buf;
|
||||
|
||||
char *salt_pos = strchr (digest_pos, ':');
|
||||
|
||||
if (input_buf[16] != data.separator) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
uint salt_len = salt_pos - digest_pos;
|
||||
|
||||
if (salt_len != 16) return (PARSER_SALT_LENGTH);
|
||||
|
||||
uint hash_len = input_len - 1 - salt_len;
|
||||
|
||||
if (hash_len != 16) return (PARSER_HASH_LENGTH);
|
||||
|
||||
salt_pos++;
|
||||
|
||||
char *salt_buf_ptr = (char *) salt->salt_buf;
|
||||
|
||||
salt_len = parse_and_store_salt (salt_buf_ptr, salt_pos, salt_len);
|
||||
|
||||
if (salt_len == UINT_MAX) return (PARSER_SALT_LENGTH);
|
||||
|
||||
salt->salt_len = salt_len;
|
||||
|
||||
uint tt;
|
||||
|
||||
salt->salt_buf_pc[0] = byte_swap_32 (salt->salt_buf[0]);
|
||||
salt->salt_buf_pc[1] = byte_swap_32 (salt->salt_buf[1]);
|
||||
|
||||
IP (salt->salt_buf_pc[0], salt->salt_buf_pc[1], tt);
|
||||
|
||||
digest[0] = hex_to_u32 ((const u8 *) &digest_pos[ 0]);
|
||||
digest[1] = hex_to_u32 ((const u8 *) &digest_pos[ 8]);
|
||||
|
||||
digest[0] = byte_swap_32 (digest[0]);
|
||||
digest[1] = byte_swap_32 (digest[1]);
|
||||
|
||||
IP (digest[0], digest[1], tt);
|
||||
|
||||
digest[0] = digest[0];
|
||||
digest[1] = digest[1];
|
||||
digest[2] = 0;
|
||||
digest[3] = 0;
|
||||
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
int lotus5_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf)
|
||||
{
|
||||
if ((input_len < DISPLAY_LEN_MIN_8600) || (input_len > DISPLAY_LEN_MAX_8600)) return (PARSER_GLOBAL_LENGTH);
|
||||
|
||||
@ -45,7 +45,7 @@ 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, 900, 1000, 1100, 1400, 1410, 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, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7100, 7200, 7300, 7400, 7500, 7600, 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, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100, 13200, 13300, 13400, 13500, 13600, 13800, 13900);
|
||||
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, 900, 1000, 1100, 1400, 1410, 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, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7100, 7200, 7300, 7400, 7500, 7600, 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, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100, 13200, 13300, 13400, 13500, 13600, 13800, 13900, 14000, 14100);
|
||||
|
||||
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);
|
||||
@ -226,7 +226,7 @@ sub verify
|
||||
$word = substr ($line, $index + 1);
|
||||
}
|
||||
# hash:salt
|
||||
elsif ($mode == 10 || $mode == 11 || $mode == 12 || $mode == 20 || $mode == 21 || $mode == 22 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 112 || $mode == 120 || $mode == 121 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1100 || $mode == 1410 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 2410 || $mode == 2611 || $mode == 2711 || $mode == 2811 || $mode == 3100 || $mode == 3610 || $mode == 3710 || $mode == 3720 || $mode == 3800 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 4900 || $mode == 5800 || $mode == 7600 || $mode == 8400 || $mode == 11000 || $mode == 12600 || $mode == 13500 || $mode == 13800 || $mode == 13900)
|
||||
elsif ($mode == 10 || $mode == 11 || $mode == 12 || $mode == 20 || $mode == 21 || $mode == 22 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 112 || $mode == 120 || $mode == 121 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1100 || $mode == 1410 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 2410 || $mode == 2611 || $mode == 2711 || $mode == 2811 || $mode == 3100 || $mode == 3610 || $mode == 3710 || $mode == 3720 || $mode == 3800 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 4900 || $mode == 5800 || $mode == 7600 || $mode == 8400 || $mode == 11000 || $mode == 12600 || $mode == 13500 || $mode == 13800 || $mode == 13900 || $mode == 14000 || $mode == 14100)
|
||||
{
|
||||
# get hash
|
||||
my $index1 = index ($line, ":");
|
||||
@ -3167,6 +3167,18 @@ sub passthrough
|
||||
{
|
||||
$tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 9));
|
||||
}
|
||||
elsif ($mode == 14000)
|
||||
{
|
||||
next if length ($word_buf) != 8;
|
||||
|
||||
$tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 16));
|
||||
}
|
||||
elsif ($mode == 14100)
|
||||
{
|
||||
next if length ($word_buf) != 24;
|
||||
|
||||
$tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 16));
|
||||
}
|
||||
else
|
||||
{
|
||||
print "ERROR: Unsupported hash type\n";
|
||||
@ -3992,6 +4004,14 @@ sub single
|
||||
}
|
||||
}
|
||||
}
|
||||
elsif ($mode == 14000)
|
||||
{
|
||||
rnd ($mode, 8, 16);
|
||||
}
|
||||
elsif ($mode == 14100)
|
||||
{
|
||||
rnd ($mode, 24, 16);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@ -7417,6 +7437,38 @@ END_CODE
|
||||
|
||||
$tmp_hash = sprintf ("%s:%s", $hash_buf, $salt_buf);
|
||||
}
|
||||
elsif ($mode == 14000)
|
||||
{
|
||||
my $salt_buf_bin = pack ("H*", $salt_buf);
|
||||
|
||||
my $cipher = new Crypt::DES $word_buf;
|
||||
|
||||
my $hash_buf = $cipher->encrypt ($salt_buf_bin);
|
||||
|
||||
$tmp_hash = sprintf ("%s:%s", unpack ("H*", $hash_buf), $salt_buf);
|
||||
}
|
||||
elsif ($mode == 14100)
|
||||
{
|
||||
my $word_buf1 = substr ($word_buf, 0, 8);
|
||||
my $word_buf2 = substr ($word_buf, 8, 8);
|
||||
my $word_buf3 = substr ($word_buf, 16, 8);
|
||||
|
||||
my $salt_buf_bin = pack ("H*", $salt_buf);
|
||||
|
||||
my $cipher1 = new Crypt::DES $word_buf1;
|
||||
|
||||
my $hash_buf1 = $cipher1->encrypt ($salt_buf_bin);
|
||||
|
||||
my $cipher2 = new Crypt::DES $word_buf2;
|
||||
|
||||
my $hash_buf2 = $cipher2->decrypt ($hash_buf1);
|
||||
|
||||
my $cipher3 = new Crypt::DES $word_buf3;
|
||||
|
||||
my $hash_buf3 = $cipher3->encrypt ($hash_buf2);
|
||||
|
||||
$tmp_hash = sprintf ("%s:%s", unpack ("H*", $hash_buf3), $salt_buf);
|
||||
}
|
||||
|
||||
return ($tmp_hash);
|
||||
}
|
||||
|
||||
@ -10,7 +10,7 @@
|
||||
|
||||
# missing hash types: 5200,6211,6221,6231,6241,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 900 1000 1100 1400 1410 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 4300 4400 4500 4700 4800 4900 5000 5100 5300 5400 5500 5600 5700 5800 6000 6100 6300 6400 6500 6600 6700 6800 6900 7100 7200 7300 7400 7500 7600 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 12100 12200 12300 12400 12600 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800"
|
||||
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 900 1000 1100 1400 1410 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 4300 4400 4500 4700 4800 4900 5000 5100 5300 5400 5500 5600 5700 5800 6000 6100 6300 6400 6500 6600 6700 6800 6900 7100 7200 7300 7400 7500 7600 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 12100 12200 12300 12400 12600 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 14000 14100"
|
||||
|
||||
#ATTACK_MODES="0 1 3 6 7"
|
||||
ATTACK_MODES="0 1 3 7"
|
||||
@ -186,6 +186,18 @@ function init()
|
||||
|
||||
fi
|
||||
|
||||
if [ "${hash_type}" -eq 14000 ]; then
|
||||
|
||||
min_len=7
|
||||
|
||||
fi
|
||||
|
||||
if [ "${hash_type}" -eq 14100 ]; then
|
||||
|
||||
min_len=23
|
||||
|
||||
fi
|
||||
|
||||
while read -u 9 pass; do
|
||||
|
||||
if [ ${i} -gt 1 ]; then
|
||||
@ -230,6 +242,18 @@ function init()
|
||||
|
||||
fi
|
||||
|
||||
if [ "${hash_type}" -eq 14000 ]; then
|
||||
|
||||
min_len=7
|
||||
|
||||
fi
|
||||
|
||||
if [ "${hash_type}" -eq 14100 ]; then
|
||||
|
||||
min_len=23
|
||||
|
||||
fi
|
||||
|
||||
# generate multiple pass/hash foreach len (2 to 8)
|
||||
if [ ${MODE} -ge 1 ]; then
|
||||
|
||||
@ -612,8 +636,12 @@ function attack_1()
|
||||
offset=7
|
||||
elif [ ${hash_type} -eq 8500 ]; then
|
||||
offset=7
|
||||
elif [ ${hash_type} -eq 14000 ]; then
|
||||
offset=7
|
||||
elif [ ${hash_type} -eq 14100 ]; then
|
||||
offset=23
|
||||
fi
|
||||
|
||||
|
||||
hash_file=${OUTD}/${hash_type}_multihash_combi.txt
|
||||
|
||||
tail -n ${offset} ${OUTD}/${hash_type}_hashes.txt > ${hash_file}
|
||||
@ -729,6 +757,20 @@ function attack_3()
|
||||
|
||||
fi
|
||||
|
||||
if [ "${hash_type}" -eq 14000 ]; then
|
||||
|
||||
mask_offset=7
|
||||
max=7
|
||||
|
||||
fi
|
||||
|
||||
if [ "${hash_type}" -eq 14100 ]; then
|
||||
|
||||
mask_offset=23
|
||||
max=23
|
||||
|
||||
fi
|
||||
|
||||
i=1
|
||||
|
||||
while read -u 9 hash; do
|
||||
@ -856,6 +898,20 @@ function attack_3()
|
||||
|
||||
fi
|
||||
|
||||
if [ "${hash_type}" -eq 14000 ]; then
|
||||
|
||||
increment_min=8
|
||||
increment_max=8
|
||||
|
||||
fi
|
||||
|
||||
if [ "${hash_type}" -eq 14100 ]; then
|
||||
|
||||
increment_min=24
|
||||
increment_max=24
|
||||
|
||||
fi
|
||||
|
||||
hash_file=${OUTD}/${hash_type}_multihash_bruteforce.txt
|
||||
|
||||
head -n $((increment_max - ${increment_min} + 1)) ${OUTD}/${hash_type}_hashes.txt > ${hash_file}
|
||||
@ -1056,6 +1112,18 @@ function attack_6()
|
||||
|
||||
fi
|
||||
|
||||
if [ "${hash_type}" -eq 14000 ]; then
|
||||
|
||||
max=6
|
||||
|
||||
fi
|
||||
|
||||
if [ "${hash_type}" -eq 14100 ]; then
|
||||
|
||||
max=6
|
||||
|
||||
fi
|
||||
|
||||
while read -u 9 hash; do
|
||||
|
||||
if [ "${i}" -gt 6 ]; then
|
||||
@ -1155,6 +1223,10 @@ function attack_6()
|
||||
max=8
|
||||
elif [ ${hash_type} -eq 8500 ]; then
|
||||
max=8
|
||||
elif [ ${hash_type} -eq 14000 ]; then
|
||||
max=5
|
||||
elif [ ${hash_type} -eq 14100 ]; then
|
||||
max=5
|
||||
fi
|
||||
|
||||
if ! contains ${hash_type} ${TIMEOUT_ALGOS}; then
|
||||
@ -1280,6 +1352,18 @@ function attack_7()
|
||||
|
||||
fi
|
||||
|
||||
if [ "${hash_type}" -eq 14000 ]; then
|
||||
|
||||
max=5
|
||||
|
||||
fi
|
||||
|
||||
if [ "${hash_type}" -eq 14100 ]; then
|
||||
|
||||
max=5
|
||||
|
||||
fi
|
||||
|
||||
i=1
|
||||
|
||||
while read -u 9 hash; do
|
||||
@ -1397,6 +1481,10 @@ function attack_7()
|
||||
max=8
|
||||
elif [ ${hash_type} -eq 8500 ]; then
|
||||
max=8
|
||||
elif [ ${hash_type} -eq 14000 ]; then
|
||||
max=5
|
||||
elif [ ${hash_type} -eq 14100 ]; then
|
||||
max=5
|
||||
fi
|
||||
|
||||
if ! contains ${hash_type} ${TIMEOUT_ALGOS}; then
|
||||
|
||||
Loading…
Reference in New Issue
Block a user